c++ - 在运行时根据 CUDA 计算能力切换主机功能

Question

我目前有一个主机功能，其中包括一个循环和各种 CUBLAS 调用。现在可以访问 CC 3.5 设备，我可以使用动态并行编写一个更高效的内核。但是，我想继续支持 CC < 3.5 设备的旧功能。我现在使用几个 gencodes 在同一个二进制文件中支持多个设备:

-gencode arch=compute_30,code=sm_30 -gencode arch=compute_35,code=sm_35

我想继续制作一个支持两种架构的二进制文件，但我想不出在主机代码中切换它的方法。 NVCC 当然无法为主机 AFAIK 上的任何内容生成编译后的代码图像。

这不好(而且非常丑陋)，因为为 CC < 3.5 构建的用户将无法使用 3.5 功能构建内核:

cudaGetDevice (&current_device);
cudaGetDeviceProperties (&current_device_properties, current_device);
if (current_device_properties.major < 3 && ... etc) {
  ...
}
else ...

__CUDACC__ 或 __CUDA_ARCH__ 在这里也没有用。

我猜这是不可能的，我将不得不在预处理器中简单地开始编译单独的二进制文件和开关架构。但是，如果有人能想到任何东西，那就太好了。

Answer 1

这取决于你的目标是什么。您似乎在这里询问两种不同的情况。

首先，如果您认为用户可能使用不支持 CC 3.5 的 nvcc编译 代码，那么您将需要对 CUDA_ARCH 使用预处理器检查来测试计算能力并防止它尝试编译不受支持的代码。

其次，如果您打算编译代码以同时包含 CC 3.5 和更低版本功能的实现，您应该使用 cudaGetDeviceProperties 检查，因为您已经注意到选择正确的主机实现。

如果您同时需要这两者，您可能需要使用看起来很像这样的实现。

cudaGetDevice (&current_device);
cudaGetDeviceProperties (&cdp, current_device);
if (cdp.major < 3 || (cdp.major >= 3 &&  cdp.minor < 5)) {
  //loop and CUBLAS
}else {
  kernel35<<<>>>();
}

同样，您的内核必须由 __CUDA_ARCH__ >= 350 保护。

#if (__CUDA_ARCH__ >= 350)
__global__ void kernel35()
{
  ...
}
#else
__global__ void kernel35()
{
  //fake stub kernel to allow non 35 compatible nvcc to build the code
}
#endif

此外，我想您已经测试过新内核的效率更高，但是如果提前知道迭代次数，动态并行几乎总是比从 CPU 正确启动慢。在我的测试中高达 40%，因此我建议在为 Kepler GPU 进行此切换之前彻底测试性能。

编辑:我突然想到，更兼容、更安全的选择是像这样表述第二部分。

__global void kernel35(){
  #if (__CUDA_ARCH__ >=350 )
  ...
  #else
  //stub
  #endif
}