Estou tentando criar uma estrutura cuda que aceite funções de usuário e encaminhe-as para o kernel, através de ponteiros de função de dispositivo. O CUDA pode trabalhar com modelos variados (-stc = c ++ 11) e até agora tudo bem.

No entanto, encontrei um problema quando o kernel chama o ponteiro da função do dispositivo. Aparentemente, o kernel roda sem problemas, mas o uso da GPU é 0%. Se eu simplesmente substituir o ponteiro de retorno de chamada pela função real, o uso da GPU será de 99%. O código aqui é muito simples e o grande intervalo de loop é simplesmente tornar as coisas mensuráveis. Eu medi o status da gpu com:

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

IIRC, a função de usuário precisa estar na mesma unidade de arquivo que o kernel (# incluído talvez) para que o nvcc seja bem-sucedido. O func_d está ali na fonte e compila e roda bem, além de não funcionar com o ponteiro de função (que é o ponto principal deste projeto).

Minha pergunta é:Por que o kernel com o ponteiro de função do dispositivo de retorno de chamada não está funcionando?

Observe que, quando imprimo os endereços de retorno de chamada e de função, eles são os mesmos, como neste exemplo de saída:

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

Outra coisa estranha é que, se alguém descomentar ocallback() ligarkernel() o uso da GPU volta a 0%, mesmo com ofunc_d() a chamada ainda está lá ... A versão func_d leva cerca de 4 segundos para ser executada, enquanto a versão de retorno de chamada não leva nada (bem, ~ 0,1s).

As especificações do sistema e o comando de compilação estão no cabeçalho do código abaixo.

Obrigado!

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