NVCC - tamanhos de bloco diferentes, dependendo do arco na hora da compilação

Question

Eu tenho um kernel que mostra o maior desempenho para diferente Tamanhos de bloco ao executar no hardware Kepler e Fermi. Eu gostaria, em tempo de compilação, para verificar o alvo de arquitetura atual e definir um THREADS_PER_BLOCK macro para i) lançar o kernel com; ii) determinar o número de blocos necessários; iii) Defina estaticamente o tamanho da memória compartilhada no kernel.

O abaixo demonstra o que estou tentando fazer. Suponha que estou direcionando o hardware GK104 e, portanto, use nvcc -arch=sm_30. Isso ainda resultará em THREADS_PER_BLOCK = 256 desde __CUDA_ARCH__ não está definido para a compilação do código do host. (Eu entendo, por exemplo esta resposta, por que não pode funcionar dessa maneira.)

#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;
}

Eu poderia verificar as propriedades do dispositivo em tempo de execução e usar a memória compartilhada dinâmica, mas gostaria de saber se isso pode ser codificado em tempo de compilação sem, por exemplo, ter que adicionar manualmente um -dfermi ou -dkepler e configuração THREADS_PER_BLOCK baseado nisso. NB: Quaisquer usuários deste código estarão compilando -o, quase Certamente para uma arquitetura, então essa não é uma opção irracional. Parece supérfluo à luz de passar o -arch= bandeira.

Answer 1

nvcc O compilador não detecta GPUs disponíveis localmente, ele sempre tem como alvo o SM 1.0 por padrão. Caso contrário, poderá introduzir um comportamento bastante confuso ao construir em diferentes sistemas.

Para compilar o dispositivo disponível, você precisa pedir ao usuário para especificar a versão SM ou executar algum código de detecção durante o tempo de construção. Estou convencido de que é mais fácil colocar o código de detecção de hardware em tempo de execução e, em seguida, configurar o lançamento do kernel conforme desejado.