重载 cuda 核函数

Posted 2023-03-23

【中文标题】重载 cuda 核函数

【英文标题】：overload cuda kernel function

【发布时间】：2021-04-25 02:05:03

【问题描述】：

在 CUDA 中使用重载的内核函数时遇到问题。

我可以理解 CUDA 可以通过参数启动重载函数。

但是，如果我想使用cudaOccupancyMaxPotentialBlockSize() 来计算最大占用率的块大小，请参阅doc。

__global__ void foo_cuda_kernel(int a)

  /*implementation 1*/


//overloaded kernel function
__global__ void foo_cuda_kernel(int a, int b)

  /*implementation 2*/


void foo_cuda()

  int min_grid_size, grid_size, block_size;
  cudaOccupancyMaxPotentialBlockSize
  (
    &min_grid_size, &block_size, 
    foo_cuda_kernel, //how does it distinguish overloaded functions?
    0, thread_num
  );
  grid_size = (thread_num + block_size - 1) / block_size;
  
  //I can understand compiler can distinguish the launched function by its arguments
  foo_cuda_kernel<<<grid_size, block_size>>>((int)1);
  cudaDeviceSynchronize();

如何让它发挥作用？ cudaOccupancyMaxPotentialBlockSize() 如何区分重载函数？

【问题讨论】：

这能回答你的问题吗？ How do I specify a pointer to an overloaded function?

不确定它是否适用于 cuda。

【参考方案1】：

如 cmets 中所述，您可以将函数强制转换为指向正确特化的指针：

auto foo_ii = static_cast<void (*)(int, int)>(&foo_cuda_kernel);
auto foo_i = static_cast<void (*)(int)>(&foo_cuda_kernel);

然后您将foo_i 或foo_ii 传递给cudaOccupancyMaxPotentialBlockSize，具体取决于您需要的函数版本。

这将起作用，因为工具链静默发出主机样板函数，这些函数包装底层运行时 API 调用以运行内核并强制内核参数类型检查。宿主编译器将这些包装器视为任何其他宿主函数（因为它们是），并自动选择匹配的版本。