为什么 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__ 函数。用 nvcc 编译时 std::sin() 是否以某种方式标记为 __device__？

Answer 1

Is std::sin() somehow marked __device__ when compiling with nvcc?

没有。在传递给 GPU 编译器的代码中，它显然被 CUDA 前端解析器替换为 sin，然后使用正常的来确保替换正确的 GPU 数学库函数。 GPU编译器看到的代码是（nvcc版本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])));
        }
    }
}

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

Why? How?

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

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

Why does std::sin() work in the CUDA kernel?

c++

cuda