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

Posted 2023-02-16

【中文标题】使用动态并行回退编译 CUDA - 多种架构/计算能力

【英文标题】：Compiling CUDA with dynamic parallelism fallback - multiple architectures/compute capability

【发布时间】：2014-01-03 03:39:38

【问题描述】：

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

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

AFAIK 你不能使用多个 -arch=sm_* 参数并且使用多个 -gencode=* 没有帮助。同样对于可分离编译，我必须使用-dlink 创建一个额外的目标文件，但是在使用compute 3.0 时不会创建这个文件（nvlink fatal : no candidate found in fatbinary 由于-lcudadevrt，我需要3.5），如何我应该处理这个吗？

【问题讨论】：

您的问题类似于this one。 AFAIK 目前没有简单的清洁解决方案。如果您解构nvcc 构建序列，则可以这样做，但我不打算详细介绍。我相信当 CUDA 6 可用时，将cudadevrt 与不尝试使用动态并行性的 pre-cc3.5 代码链接时将不再抛出错误，然后这个问题将很容易解决。 CUDA 6 应该很快就会推出。

【参考方案1】：

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

这是我的简单测试：

$ 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 设备。

【参考方案2】：

从 CUDA 5.5 开始，我认为没有办法使用运行时 API 来做到这一点。

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