具有函数指针和可变参数模板的 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 no 回调和 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的地址匹配呢？”因为您在主机代码中使用 both 地址，并且这两个地址都是垃圾。为了让自己相信这一点，请在内核的最后添加一行类似这样的内容：

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

   您会看到它打印出的内容与主机上打印的内容不同。

3. “设备利用率如何？”一旦设备遇到错误，内核就会终止，利用率会变为零。希望这将为您解释此声明：“另一件奇怪的事情是，如果取消了 kernel() 中的 callback() 调用，那么 GPU 使用率将回到 0%，即使 func_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
$