CUDA-Kernel mit Funktionszeiger und variadischen Vorlagen

Ich versuche, ein cuda-Framework zu entwerfen, das Benutzerfunktionen akzeptiert und über Zeiger auf Gerätefunktionen an den Kernel weiterleitet. CUDA kann mit verschiedenen Vorlagen (-stc = c ++ 11) arbeiten und ist bisher so gut.

Ich bin jedoch auf ein Problem gestoßen, als der Kernel den Gerätefunktionszeiger aufruft. Anscheinend läuft der Kernel ohne Probleme, aber die GPU-Auslastung beträgt 0%. Wenn ich einfach den Rückrufzeiger durch die eigentliche Funktion ersetze, beträgt die GPU-Auslastung 99%. Der Code hier ist sehr einfach und der große Schleifenbereich dient einfach dazu, Dinge messbar zu machen. Ich habe den gpu status gemessen mit:

nvidia-smi --query-gpu=utilization.gpu,utilization.mory,memory.used --format=csv -lms 100 -f out.txt

IIRC, die Benutzerfunktion muss sich in derselben Dateieinheit wie der Kernel befinden (# möglicherweise eingeschlossen), damit nvcc erfolgreich ist. Die Funktion func_d befindet sich direkt in der Quelle und wird ordnungsgemäß kompiliert und ausgeführt, abgesehen davon, dass sie nicht mit dem Funktionszeiger funktioniert (worauf es in diesem Entwurf ankommt).

Meine Frage ist:Warum funktioniert der Kernel mit dem Funktionszeiger des Rückrufgeräts nicht?

Beachten Sie, dass, wenn ich die Callback- und die func_d-Adresse nicht drucke, diese dieselben sind wie in der folgenden Beispielausgabe:

size of Args = 1
callback() address = 4024b0
func_d()   address = 4024b0

Eine andere seltsame Sache ist, wenn man das @ auskommentiecallback() anrufenkernel() dann ist die GPU-Auslastung auch mit dem @ wieder auf func_d() call noch drin ... Die Ausführung der func_d-Version dauert ca. 4 Sekunden, während die Callback-Version nichts benötigt (~ 0.1sec).

Systemspezifikationen und der Kompilierungsbefehl befinden sich im Kopf des folgenden Codes.

Vielen Dank

// compiled with:
// nvcc -g -G -O0 -std=c++11 -arch=sm_20 -x cu sample.cpp
//
// Nvidia Quadro 6000 (compute capability 2.0)
// CUDA 6.5 (V6.5.12),
// Arch Linux, Nvidia driver 343.22-4, gcc 4.9.1
// Nov, 2014


#include <stdio.h>

__device__
void func_d(double* vol)
{
*vol += 5.4321f;
}


// CUDA kernel function
template <typename... Types>
__global__ void kernel( void (*callback)(Types*...) )
{
double val0 = 1.2345f;

//  // does not use gpu (0% gpu utilization)
//  for ( int i = 0; i < 1000000; i++ ) {
//  callback( &val0 );
//  }

// uses gpu (99% gpu utilization)
for ( int i = 0; i < 10000000; i++ ) {
func_d( &val0 );
}
}


// host function
template <typename... Types>
void host_func( void (*callback)(Types*...) )
{
// get user kernel number of arguments.
constexpr int I = sizeof...(Types);
printf("size of Args = %d\n",I);

printf("callback() address = %x\n",callback);
printf("func_d()   address = %x\n",func_d);

dim3 nblocks = 100;
int nthread = 100;
kernel<Types...><<<nblocks,nthread>>>( callback );
}


__host__
int main(int argc, char** argv)
{
host_func(func_d);
}