同一库中的动态并行性和CC 2.0代码

Question

在我的库中，我需要支持计算能力2.0及更高的设备。对于CC 3.5设备，我已经实现了利用动态并行的优化内核。当指定小于“compute_35，sm_35”时，nvcc编译器似乎不支持DP（我收到编译器/链接器错误）。我的问题是在这种情况下支持多个内核版本的最佳方法是什么？拥有多个DLL并在运行时在它们之间进行选择将起作用，但我想知道是否有更好的方法。

更新：我已成功使用#if __CUDA_ARCH__

1.




Answer 1

我相信CUDA 6已经解决了这个问题。

特别是，与指定 -lcudadevrt 库并为不需要动态并行的代码抛出链接错误相关的编译问题已被消除/删除。

这是我的简单测试：

$ cat t264.cu
#include <stdio.h>

__global__ void kernel1(){
  printf("Hello from DP Kernel\n");
}

__global__ void kernel2(){

#if __CUDA_ARCH__ >= 350
  kernel1<<<1,1>>>();
#else
  printf("Hello from non-DP Kernel\n");
#endif
}

int main(){

  kernel2<<<1,1>>>();
  cudaDeviceSynchronize();
  return 0;
}
$ nvcc -O3 -gencode arch=compute_20,code=sm_20 -gencode arch=compute_35,code=sm_35 -rdc=true -o t264 t264.cu -lcudadevrt
$ CUDA_VISIBLE_DEVICES="0" ./t264
Hello from non-DP Kernel
$ CUDA_VISIBLE_DEVICES="1" ./t264
Hello from DP Kernel
$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2013 NVIDIA Corporation
Built on Sat_Jan_25_17:33:19_PST_2014
Cuda compilation tools, release 6.0, V6.0.1
$

在我的例子中，设备0是Quadro5000，cc 2.0设备，设备1是GeForce GT 640，cc 3.5设备。