Ядро 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);
}