NVCC: diferentes tamaños de bloques dependiendo del arco en el tiempo de compilación

Question

Tengo un núcleo que muestra el mayor rendimiento para diferente Tamaños de bloque cuando se ejecuta en Hardware Kepler y Fermi. Me gustaría, en tiempo de compilación, verificar el objetivo de arquitectura actual y definir un THREADS_PER_BLOCK macro a i) lanzar el núcleo con; ii) determinar el número de bloques necesarios; iii) Establezca estáticamente el tamaño de memoria compartida en el núcleo.

El siguiente demuestra lo que estoy intentando hacer. Supongamos que estoy apuntando al hardware GK104 y, por lo tanto, uso nvcc -arch=sm_30. Esto aún dará como resultado THREADS_PER_BLOCK = 256 ya que __CUDA_ARCH__ no está definido para la compilación del código de host. (Entiendo, de EG esta respuesta, por qué no puede funcionar de esta manera).

#if __CUDA_ARCH__ >= 300
#define THREADS_PER_BLOCK 512
#else
#define THREADS_PER_BLOCK 256
#endif

__global__ void some_kernel(int* a, int* b) {
    __shared__ sm_data[THREADS_PER_BLOCK];
    // Do something.
}

int main(void) {
    // Initialize data.
    // Calculate blocks based on THREADS_PER_BLOCK, problem size and some max.
    some_kernel<<blocks, THREADS_PER_BLOCK>>>(d_a, d_b)
    return 0;
}

Podría verificar las propiedades del dispositivo en el tiempo de ejecución y usar memoria compartida dinámica, pero me gustaría saber si esto puede estar codificado en el tiempo de compilación sin tener que agregar manualmente un -dfermi o -dkepler y configurar THREADS_PER_BLOCK basado en eso. NB: Cualquier usuario de este código lo compilará ellos mismos, casi Ciertamente para una arquitectura, por lo que esta no es una opción irrazonable. Simplemente parece superfluo a la luz de pasar el -arch= bandera.

Answer 1

nvcc El compilador no detecta GPU localmente disponible, siempre se dirige a SM 1.0 por defecto. De lo contrario, podría introducir un comportamiento bastante confuso al construir en diferentes sistemas.

Para compilar el dispositivo disponible, debe pedirle al usuario que especifique la versión SM o ejecute algún código de detección durante el tiempo de compilación. Estoy convencido de que es más fácil poner el código de detección de hardware en tiempo de ejecución y luego configurar el lanzamiento de su núcleo según lo desee.