为什么 std::sin() 在 CUDA 内核中工作？

Question

下面的代码编译（使用nvcc test.cu -o test）并运行没有错误，这意味着它std::sin()在设备上工作：

#include <cmath>
#include <vector>
#include <cassert>
#include <numeric>

__global__ void map_sin(double* in, double* out, int n) {
  const int i = blockIdx.x * 512 + threadIdx.x;
  if (i < n) {
    out[i] = std::sin(in[i]);
  }
}

int main() {
  const int n = 1024;
  std::vector<double> in(n), out(n);
  std::iota(in.begin(), in.end(), 1.);

  double *in_, *out_;
  cudaMalloc(reinterpret_cast<void**>(&in_), n * sizeof(double));
  cudaMemcpy(in_, in.data(), n * sizeof(double), cudaMemcpyHostToDevice);
  cudaMalloc(reinterpret_cast<void**>(&out_), n * sizeof(double));

  map_sin<<<n / 512, 512>>>(in_, out_, n);

  cudaMemcpy(out.data(), out_, n * sizeof(double), cudaMemcpyDeviceToHost);
  cudaFree(in_);
  cudaFree(out_);

  for (int i = 0; i != 10; ++i) {
    assert(std::abs(out[i] - std::sin(in[i])) < 1e-3);
  }
}

为什么？如何？根据这个答案，CUDA 内核应该只能调用__device__函数。被std::sin()莫名其妙地标记__device__与编译时nvcc？

Answer 1

使用 nvcc 编译时是否std::sin()以某种方式标记__device__？

不是。sin在传递给GPU编译器的代码中，它显然被CUDA前端解析器替换，然后使用正常的重载机制来确保正确的GPU数学库函数被替换。GPU编译器看到的代码是（nvcc version 11.1.74）：

__global__ __var_used__ void _Z7map_sinPdS_i( double *in, double *out, int n)
{
    {
        int __cuda_local_var_31644_13_non_const_i;
        __cuda_local_var_31644_13_non_const_i = ((int)(((blockIdx.x) * 512U) + (threadIdx.x)));
        if (__cuda_local_var_31644_13_non_const_i < n)
        {
        (out[__cuda_local_var_31644_13_non_const_i]) = (sin((in[__cuda_local_var_31644_13_non_const_i])));
        }
    }
}

如您所见，没有命名空间。

为什么？如何？

它没有记录在案，我没有关于实现细节的特殊内幕信息，但我猜测对于<cmath> （其中CUDA 数学库具有 C++11 标准中定义的所有内容的实现）的内容，命名空间只是脱光，一切正常。对于其他 C++ 标准库函数，这不会发生，并且您将遇到您期望的那种编译失败。