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

Question

我目前有一个主机函数，其中包括一个循环和各种 CUBLAS 调用。现在可以访问 CC 3.5 设备，我可以使用动态并行编写一个更高效的内核。但是，我想继续支持 CC

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

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

这不好（而且非常丑陋），因为为 CC

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
}