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

user_905845 0
Viewed 691

The following code compiles (with nvcc test.cu -o test) and runs without error, meaning that std::sin() does work on the device:

#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);
  }
}

Why? How? According to this answer, CUDA kernels are supposed to be able to call only __device__ functions. Is std::sin() somehow marked __device__ when compiling with nvcc?

cuda c++
1 Answers

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

No. It is apparently replaced with sin by the CUDA front end parser in the code which is passed to the GPU compiler, and then the normal overload mechanism is used to ensure the correct GPU math library function is substituted. The code which the GPU compiler sees is (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])));
        }
    }
}

As you can see, there are no namespaces.

Why? How?

It isn't documented and I have no special insider information on implementation details, but I am guessing that for the contents of <cmath> (where the CUDA math library has an implementation of everything defined in the C++11 standard), the namespace is just stripped off and everything works. For other C++ standard library functions, that can't happen and you will get a compile failure of the sort you expect.

user_6818650
Related