使用动态并行回退编译 CUDA - 多种架构/计算能力

Question

在一个应用程序中，我有一堆CUDA内核。有些使用动态并行，有些则不使用。为了在不支持的情况下提供回退选项，或者只是允许应用程序继续使用减少/部分可用的功能，我该如何进行编译？

目前，在不需要计算 3.5 的 670（最大 sm_30）上运行用 -arch=sm_35 编译的内核时，我遇到了invalid device function。

AFAIK 您不能使用多个-arch=sm_*参数，并且使用多个-gencode=*也无济于事。同样对于可分离的编译，我不得不使用 -dlink 创建一个额外的对象文件，但是在使用 compute 3.0 时不会创建该文件（由于 3.5 需要的 -lcudadevrt nvlink fatal : no candidate found in fatbinary），我应该如何处理这个问题？

Answer 1

我相信这个问题现在已经在 CUDA 6 中得到解决。

这是我的简单测试：

$ cat t264.cu
#include <stdio.h>
__global__ void kernel1(){
  printf("Hello from DP Kerneln");
}
__global__ void kernel2(){
#if __CUDA_ARCH__ >= 350
  kernel1<<<1,1>>>();
#else
  printf("Hello from non-DP Kerneln");
#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 设备。

Answer 2

我不相信有办法使用 CUDA 5.5 的运行时 API 来做到这一点。

我能想到的解决这个问题的唯一方法是使用驱动程序 API 执行您自己的架构选择并在运行时从不同的 cubin 文件加载代码。API 可以安全地混合，因此只需要使用驱动程序 API 完成上下文建立-设备选择-模块加载阶段。之后，您可以使用运行时 API - 内核启动需要一点自制的语法糖，否则不需要在其他运行时 API 代码中进行代码更改。