Ядро CUDA с указателем на функцию и переменными шаблонами

Я пытаюсь разработать фреймворк cuda, который бы принимал пользовательские функции и передавал их ядру через указатели функций устройства. CUDA может работать с переменными шаблонами (-stc = c ++ 11) и пока хорошо.

Однако я сталкиваюсь с проблемой, когда ядро ​​вызывает указатель на функцию устройства. Очевидно, ядро ​​работает без проблем, но использование графического процессора составляет 0%. Если я просто заменю указатель обратного вызова на фактическую функцию, то использование графического процессора составляет 99%. Код здесь очень прост, а большой диапазон циклов просто делает вещи измеримыми. Я измерил статус GPU с помощью:

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

IIRC, пользовательская функция должна быть в том же файловом модуле, что и ядро ​​(возможно, #included) для успешного выполнения nvcc. Func_d прямо в исходном коде, он компилируется и работает нормально, кроме того, что не работает с указателем на функцию (что и является целью этого проекта).

Мой вопрос:Почему не работает ядро ​​с указателем функции устройства обратного вызова?

Обратите внимание, что когда я печатаю адреса callback и func_d, они такие же, как в этом примере вывода:

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

Еще одна странная вещь, если кто-то раскомментируетcallback() вызыватьkernel() затем использование графического процессора возвращается к 0%, даже еслиfunc_d() вызов все еще там ... Версия func_d занимает около 4 секунд, в то время как версия обратного вызова не занимает ничего (ну, ~ 0,1 сек).

Системные спецификации и команда компиляции находятся в главе кода ниже.

Спасибо!

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

Ответы на вопрос(1)

Ваш ответ на вопрос