0

我目前有一个主机函数，其中包括一个循环和各种 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__ 在这里也没有用。

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

4

1 回答 1

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
}
于 2013-06-14T15:16:06.293 回答