Estoy tratando de diseñar un marco cuda que acepte las funciones del usuario y las reenvíe al kernel, a través de punteros de funciones del dispositivo. CUDA puede trabajar con plantillas variadas (-stc = c ++ 11) y hasta ahora todo bien.

Sin embargo, me topé con un problema cuando el núcleo llama al puntero de la función del dispositivo. Aparentemente, el núcleo se ejecuta sin problemas, pero el uso de la GPU es del 0%. Si simplemente reemplazo el puntero de devolución de llamada con la función real, el uso de GPU es del 99%. El código aquí es muy simple y el amplio rango de bucle es simplemente para hacer que las cosas sean medibles. Medí el estado de gpu con:

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

IIRC, la función de usuario debe estar en la misma unidad de archivo que el kernel (#incluido quizás) para que nvcc tenga éxito. El func_d está justo allí en la fuente y se compila y funciona bien, además de no funcionar con el puntero de función (que es el punto central de este diseño).

Mi pregunta es:¿Por qué el kernel con el puntero de función del dispositivo de devolución de llamada no funciona?

Tenga en cuenta que, cuando imprimo las direcciones de devolución de llamada y func_d, son las mismas, como en esta salida de muestra:

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

Otra cosa extraña es, si uno comenta elcallback() llamarkernel() entonces el uso de la GPU vuelve al 0%, incluso con elfunc_d() la llamada sigue ahí ... La versión de func_d tarda unos 4 segundos en ejecutarse, mientras que la versión de devolución de llamada no toma nada (bueno, ~ 0.1 segundos).

Las especificaciones del sistema y el comando de compilación están en el encabezado del código a continuación.

¡Gracias!

// 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);
}