取决于 nvcc 标志的无效设备符号错误

Question

玩具节目：

#include <iostream>
#include <vector>

// Matrix side size (they are square).
const int N = 3;
const int num_mats = 14;

// Rotation matrices.
__constant__ float rot_mats_device[num_mats*N*N];

int main() {
  std::vector<float> rot_mats_host(num_mats*N*N);
  for (int i = 0; i < rot_mats_host.size(); i++)
    rot_mats_host[i] = i;

  auto errMemcpyToSymbol = cudaMemcpyToSymbol(rot_mats_device,
                                              rot_mats_host.data(),
                                              sizeof(rot_mats_device));

  if (errMemcpyToSymbol != cudaSuccess) {
    std::cout << "MemcpyToSymbol error: " <<
      cudaGetErrorString(errMemcpyToSymbol) << std::endl;
  }
}

编译为

nvcc -arch=sm_52 -std=c++11 cuda_invalid_symbol_error.cu -o cuda_invalid_symbol_error

在运行时不会给出任何错误。然而，

nvcc -gencode arch=compute_52,code=sm_52 -std=c++11 cuda_invalid_symbol_error.cu -o cuda_invalid_symbol_error

它将失败并显示消息 MemcpyToSymbol error: invalid device symbol。

为什么后面的编译指令会报运行时错误？

规格：cuda 8.0，Ubuntu 16.04，GeForce GTX 1060（我知道这张卡的cc是6.1）。

Answer 1

Why do the latter instructions for compilation give the runtime error?

-arch=sm_xx 是 shorthand 对于：

-gencode arch=compute_xx,code=sm_xx -gencode arch=compute_xx,code=compute_xx

在你的例子中，xx 是 52，这个命令嵌入了 cc 5.2 PTX 代码（第二个 gencode 实例）和 cc 5.2 SASS 代码（第一个 gencode实例）。 cc 5.2 的 SASS 代码不会在您的 cc6.1 设备上运行，因此运行time JIT 编译 cc 5.2 PTX 代码以创建与您的 cc 6.1 架构兼容的对象.万事如意，万事如意。

当您改为编译时：

nvcc -gencode arch=compute_52,code=sm_52 ...

您正在从编译对象中省略 PTX 代码。只有 cc 5.2 SASS 代码存在。此代码不会运行在你的cc6.1设备上，运行时间没有其他选项，所以运行时会出现NO_BINARY_FOR_GPU的"hidden"错误]time 尝试为您的程序加载 GPU 图像。由于没有加载图像，因此没有设备符号 present/usable。由于它不是 present/usable，当您尝试使用 CUDA 运行time API 引用它时，您会得到 invalid device symbol 错误。

如果您在此之前执行了另一个 CUDA 运行time API 调用，强制对 CUDA 运行time 进行足够或等效级别的初始化（并检查返回的错误代码），您会收到 NO_BINARY_FOR_GPU 错误或类似错误。当然，例如，如果您尝试启动 GPU 内核，您会收到该错误。可能还有其他 CUDA 运行time API 调用会强制执行等效或足够级别的延迟初始化，但我没有这些列表。

取决于 nvcc 标志的无效设备符号错误

Invalid device symbol error depending on nvcc flags

cuda

nvcc