CUDA: L'oubli de la configuration de lancement du noyau n'entraîne pas d'avertissement ou d'erreur du compilateur NVCC

Question

Lorsque j'essaie d'appeler un noyau CUDA (fonction __global__) à l'aide d'un pointeur de fonction, tout semble fonctionner correctement. Cependant, si j'oublie de fournir une configuration de lancement lors de l'appel du noyau, NVCC ne produira pas d'erreur ou d'avertissement, mais le programme se compilera et se bloquera si je tente de l'exécuter.

__global__ void bar(float x) { printf("foo: %f\n", x); } 

typedef void(*FuncPtr)(float); 

void invoker(FuncPtr func) 
{ 
    func<<<1, 1>>>(1.0); 
} 

invoker(bar); 
cudaDeviceSynchronize(); 

Compilez et exécutez ce qui précède. Tout ira bien. Ensuite, supprimez la configuration de lancement du noyau (c.-à-d. < < < 1, 1 >>>). Le code compilera très bien, mais il va planter lorsque vous essayez de l'exécuter.

Une idée de ce qui se passe? Est-ce un bug, ou je ne suis pas censé passer les pointeurs des fonctions __global__?

la version CUDA: 8.0

version OS: Debian (test repo) GPU: NVIDIA GeForce 750M

Answer 1

Si nous prenons une version légèrement plus complexe de votre repro, et regardez le code émis par le Cuda toolchain-extrémité avant, il devient possible de voir ce qui se passe:

#include <cstdio> 

__global__ void bar_func(float x) { printf("foo: %f\n", x); } 
typedef void(*FuncPtr)(float); 

void invoker(FuncPtr passed_func) 
{ 
#ifdef NVCC_FAILS_HERE 
    bar_func(1.0); 
#endif 
    bar_func<<<1,1>>>(1.0); 
    passed_func(1.0); 
    passed_func<<<1,1>>>(2.0); 
} 

Alors compilons il deux façons:

$ nvcc -arch=sm_52 -c -DNVCC_FAILS_HERE invoker.cu 
invoker.cu(10): error: a __global__ function call must be configured 

En d'autres termes, le frontal peut détecter que bar_func est une fonction globale et nécessite des paramètres de lancement. Une autre tentative:

$ nvcc -arch=sm_52 -c -keep invoker.cu 

Comme vous le constatez, cela ne génère aucune erreur de compilation. Regardons ce qui est arrivé:

void bar_func(float x) ; 
# 5 "invoker.cu" 
typedef void (*FuncPtr)(float); 
# 7 "invoker.cu" 
void invoker(FuncPtr passed_func) 
# 8 "invoker.cu" 
{ 
# 12 "invoker.cu" 
(cudaConfigureCall(1, 1)) ? (void)0 : (bar_func)((1.0)); 
# 13 "invoker.cu" 
passed_func((2.0)); 
# 14 "invoker.cu" 
(cudaConfigureCall(1, 1)) ? (void)0 : passed_func((3.0)); 
# 15 "invoker.cu" 
} 

La syntaxe d'invocation standard du noyau se <<<>>> étendu à un appel en ligne à cudaConfigureCall, puis une fonction wrapper hôte est appelé. L'enveloppe hôte a les API internes nécessaires pour lancer le noyau:

void bar_func(float __cuda_0) 
# 3 "invoker.cu" 
{__device_stub__Z8bar_funcf(__cuda_0); } 

void __device_stub__Z8bar_funcf(float __par0) 
{ 
    if (cudaSetupArgument((void *)(char *)&__par0, sizeof(__par0), (size_t)0UL) != cudaSuccess) return; 
    { volatile static char *__f __attribute__((unused)); __f = ((char *)((void (*)(float))bar_func)); 
     (void)cudaLaunch(((char *)((void (*)(float))bar_func))); 
    }; 
} 

Ainsi, le talon ne gère que les arguments et lance le noyau via cudaLaunch. Il ne gère pas la configuration de lancement

La raison sous-jacente de l'incident (en réalité une erreur d'API d'exécution non détectée) est que le lancement du noyau se produit sans configuration préalable. Évidemment, cela arrive parce que le frontal CUDA (et C++ d'ailleurs) ne peut pas faire d'introspection de pointeur au moment de la compilation et détecter que votre pointeur de fonction est une fonction de remplacement pour appeler un noyau.

Je pense que la seule façon de décrire cela est une "limitation" de l'API d'exécution et du compilateur. Je ne dirais pas que ce que vous faites est faux, mais j'utiliserais probablement l'API du pilote et gérerais moi-même le noyau de façon explicite dans une telle situation.