CUDA内核，带有函数指针和可变模板

CUDA kernel with function pointer and variadic templates

本文关键字：指针 函数 内核 CUDA      更新时间：2023-10-16

我正在尝试设计一个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，用户函数需要与内核位于同一文件单元中（可能包括#），才能成功nvcc。func_d就在源代码中，它编译和运行得很好，除了不使用函数指针之外（这是本设计的全部要点）。

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

注意，当我打印回调和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 = %dn",I);
printf("callback() address = %xn",callback);
printf("func_d()   address = %xn",func_d);
dim3 nblocks = 100;
int nthread = 100;
kernel<Types...><<<nblocks,nthread>>>( callback );
}

__host__
int main(int argc, char** argv)
{
host_func(func_d);
}

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

可能有几个问题需要解决。但最简单的答案是，在主机代码中获取设备实体的地址是非法的。设备变量和设备函数都是如此。现在，可以获取这些实体的地址。但是这个地址是垃圾。它在主机或设备上都不可用。如果你无论如何都试图使用它们，你会在设备上得到未定义的行为，这通常会使你的内核停止。

在主机代码中可以看到主机地址。在设备代码中可以看到设备地址。任何其他行为都需要API干预。

  1. 您似乎在使用nvidia-smi利用率查询来衡量事情是否正常运行。我建议您进行适当的cuda错误检查，并且您可能希望使用cuda-memcheck运行您的代码。

  2. "为什么func_d的地址与callback的地址匹配？"因为您在主机代码中同时使用两个地址，而这两个地址都是垃圾。为了让自己相信这一点，在内核的最后添加一行类似的内容：

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

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

  3. "设备利用率如何？"一旦设备遇到错误，内核就会终止，利用率就会变为零。希望这能为你解释这句话："另一件奇怪的事情是，如果在kernel（）中取消对callback（）调用的注释，那么GPU的使用率就会回到0%，即使func_d（）调用仍然存在…"

  4. "我怎么能解决这个问题？"我不知道有什么好方法可以解决这个问题。如果在编译时已知的CUDA函数数量有限，并且希望用户能够从中进行选择，那么适当的做法可能只是创建一个适当的索引，并使用它来选择函数。如果你真的想，你可以运行一个初步的/设置内核，它将获取你关心的函数的地址，然后你可以将这些地址传递回主机代码，并在随后的内核调用中使用它们作为参数，这应该允许你的机制工作。但我看不出它是如何防止需要通过编译时已知的一组预定义函数进行索引的。如果您的方向是希望用户能够在运行时提供用户定义的函数，我想您会发现这在使用CUDA运行时API时非常困难（我怀疑这在未来可能会改变）。)我提供了一个相当扭曲的机制来尝试在这里做到这一点（阅读整个问答；Talonmy在那里的回答也很有信息）。另一方面，如果您愿意使用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 = %fn", *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 = %xn",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 = %dn",I);
  printf("callback() address = %xn",callback);
  printf("func_d()   address = %xn",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
$
相关文章：