时间戳

首页 > 解决方案 > 为什么 std::sin() 在 CUDA 内核中工作？

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

问题描述

以下代码编译（使用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

标签： c++cuda

解决方案

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

不。它显然被sin传递给GPU编译器的代码中的CUDA前端解析器替换，然后使用正常的重载机制来确保替换正确的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])));
        }
    }
}

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

为什么？如何？

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

推荐阅读