NVCC - différentes tailles de blocs en fonction de l'arc au moment de la compilation

Question

J'ai un noyau qui montre des performances les plus élevées pour différent Bloquer les tailles lors de l'exécution sur Kepler et Fermi Hardware. Je voudrais, au moment de la compilation, vérifier la cible d'architecture actuelle et définir un THREADS_PER_BLOCK macro pour i) lancer le noyau avec; ii) Déterminer le nombre de blocs nécessaires; iii) Définissez statiquement la taille de la mémoire partagée dans le noyau.

Ce qui est ci-dessous montre ce que j'essaie de faire. Supposons que je cible le matériel GK104, et donc j'utilise nvcc -arch=sm_30. Cela entraînera toujours THREADS_PER_BLOCK = 256 puisque __CUDA_ARCH__ n'est pas défini pour la compilation de code hôte. (Je comprends, par exemple Cette réponse, pourquoi ça ne peut pas fonctionner de cette façon.)

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

Je pourrais vérifier les propriétés des appareils au moment de l'exécution et utiliser la mémoire partagée dynamique, mais je voudrais savoir si cela peut être codé austère au moment de la compilation sans avoir à ajouter manuellement un -dfermi ou -dkepler et à régler THREADS_PER_BLOCK basé sur cela. NB: Tous les utilisateurs de ce code le compileront eux-mêmes, presque Certainement pour une architecture, ce n'est donc pas une option déraisonnable. Il semble juste superflu à la lumière de passer le -arch= drapeau.

Answer 1

nvcc Le compilateur ne détecte pas les GPU disponibles localement, il cible toujours SM 1.0 par défaut. Sinon, il pourrait introduire un comportement assez déroutant lors de la construction de différents systèmes.

Pour compiler pour l'appareil disponible, vous devez demander à l'utilisateur de spécifier la version SM ou d'exécuter un code de détection pendant la durée de construction. Je suis convaincu qu'il est plus facile de mettre le code de détection matérielle dans l'exécution, puis de configurer votre lancement de noyau comme vous le souhaitez.