CUDA内核与函数指针和可变参数模板

Question

我想设计一个cuda框架，它将接受用户函数并通过设备函数指针将它们转发给内核。 CUDA可以使用可变参数模板（-stc = C++ 11），并且非常好。

但是，当内核调用设备函数指针时出现问题。显然内核运行没有问题，但GPU使用率为0％。如果我简单地用实际函数替换回调指针，那么GPU使用率为99％。这里的代码非常简单，大的循环范围仅仅是为了使事情可以测量。我测与GPU状态：

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

IIRC，用户功能需要在同一个文件为单位的内核（可能是执行#included）以NVCC成功。 func_d就在源代码中，并且它编译并运行良好，除了不使用函数指针（这是本设计中的全部要点）。

我的问题是： 为什么带回调设备函数指针的内核不工作？

需要注意的是，当我printf的诺斯回调和func_d地址，它们是相同的，因为在此示例输出：

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

另一个奇怪的是，如果取消注释在kernel()的callback()呼叫，则GPU使用率回到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); 
} 

Answer 1

我的问题是：为什么用装置的回调函数指针内核不工作？

可能有几个问题需要解决。但最简单的答案是因为在主机代码中取得设备实体的地址是非法的。对于设备变量以及设备功能来说，这是正确的。现在，你可以取这些实体的地址。但地址是垃圾。它不能在主机或设备上使用。如果你尝试使用它们，你会在设备上产生未定义的行为，这通常会让你的内核停顿。

在主机代码中可能会出现主机地址。设备地址可以在设备代码中观察到。任何其他行为都需要API干预。

1. 你似乎是使用nvidia-smi利用查询作为东西是否被正确运行的措施。我建议您改为使用proper cuda error checking，并且您也可以使用cuda-memcheck运行您的代码。

2. “为什么func_d的地址与callback的地址匹配？“因为你正在服用都地址在主机代码，和地址都是垃圾说服自己这一点，在你的内核的最后添加一行是这样的：

   if ((!threadIdx.x)&&(!blockIdx.x)) printf("in-kernel func_d() address = %x\n",func_d); 
   

   ，你会看到它打印出与主机上打印的东西不同的内容

3. “设备使用情况如何？”只要设备遇到错误，内核就会终止并且利用率为零。为你解释这句话：“另一个奇怪的是，如果在kernel（）中取消注释callback（）的调用，那么GPU的使用率会回到0％，即使使用fu nc_d（）调用还在里面......“

4. ‘我怎样才能解决这个问题？’我不知道一个伟大的方式来解决这个问题。如果您在编译时已知有限数量的CUDA函数，并希望用户能够从中进行选择，那么恰当的事情可能只是创建适当的索引，然后使用它来选择函数。如果你真的想要的话，你可以运行一个初步的/设置的内核，它将获取你关心的函数的地址，然后你可以将这些地址传递回主机代码，并在随后的内核调用中用它们作为参数，应该允许你的机制工作。但我不明白它是如何防止需要通过编译时已知的一组预定义函数进行索引的。如果你正在前进的方向是，你希望用户能够在运行时提供用户自定义函数我想你会觉得这很困难的时刻与CUDA运行时API做（我怀疑这很可能在未来改变），我提供了一个相当扭曲的机制，试图做到这一点here（读取整个问答; talonmies回答有内容，也）。另一方面，如果您愿意使用CUDA驱动程序API，那么它应该是可能的，尽管有些参与，因为这正是PyCUDA中非常优雅的方式。

5. 在未来，请缩进代码。

下面是一个完整的工作示例，演示了上述几个想法。特别是，我显示在一个相当粗糙的方式，即func_d地址可在设备代码中存在时，则传回主机，然后用作未来内核参数来成功地选择/调用该设备的功能。

$ cat t595.cu 
// 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) 
{ 
    if ((!threadIdx.x) && (!blockIdx.x)) printf("value = %f\n", *vol); 
    *vol += 5.4321f; 
} 

template <typename... Types> 
__global__ void setup_kernel(void (**my_callback)(Types*...)){ 
    *my_callback = func_d; 
} 

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

    val0 = 0.0f; 
// uses gpu (99% gpu utilization) 
// for (int i = 0; i < 10000000; i++) { 
    func_d(&val0); 
// } 
    if ((!threadIdx.x)&&(!blockIdx.x)) printf("in-kernel func_d() address = %x\n",func_d); 
} 


// 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; 
    unsigned long long *d_callback, h_callback; 
    cudaMalloc(&d_callback, sizeof(unsigned long long)); 
    setup_kernel<<<1,1>>>((void (**)(Types*...))d_callback); 
    cudaMemcpy(&h_callback, d_callback, sizeof(unsigned long long), cudaMemcpyDeviceToHost); 
    kernel<Types...><<<nblocks,nthread>>>((void (*)(Types*...))h_callback); 
    cudaDeviceSynchronize(); 
} 


__host__ 
int main(int argc, char** argv) 
{ 
    host_func(func_d); 
} 
$ nvcc -std=c++11 -arch=sm_20 -o t595 t595.cu 
$ cuda-memcheck ./t595 
========= CUDA-MEMCHECK 
size of Args = 1 
callback() address = 4025dd 
func_d() address = 4025dd 
value = 1.234500 
value = 0.000000 
in-kernel func_d() address = 4 
========= ERROR SUMMARY: 0 errors 
$