CUDA 是否实现了 std::complex 运算符 =？

Question

我修改了int版本的vector add到两个复数vector相加，下面的代码可以工作，但是我很困惑：

#include <stdio.h>
#include <complex>

#define N (2048*2048)
#define THREADS_PER_BLOCK 512

__global__ void add(std::complex<double> *a, std::complex<double> *b, std::complex<double> *c)
{
  int index = threadIdx.x + blockIdx.x * blockDim.x;
  // c[index] = a[index] + b[index];
  // c[index] = a[index].real();
  c[index] = a[index];
}


int main()
{
    // host side
    std::complex<double> *a;
    std::complex<double> *b;
    std::complex<double> *c;

    // device side
    std::complex<double> *d_a;
    std::complex<double> *d_b;
    std::complex<double> *d_c;

    int size = N * sizeof(std::complex<double>);

/* allocate space for device copies of a, b, c */

  cudaMalloc( (void **) &d_a, size );
  cudaMalloc( (void **) &d_b, size );
  cudaMalloc( (void **) &d_c, size );

/* allocate space for host copies of a, b, c and setup input values */

  a = (std::complex<double>*)malloc( size );
  b = (std::complex<double>*)malloc( size );
  c = (std::complex<double>*)malloc( size );

  for( int i = 0; i < N; i++ )
  {
    a[i] = b[i] = i;
    c[i] = 0;
  }


  cudaMemcpy( d_a, a, size, cudaMemcpyHostToDevice );
  cudaMemcpy( d_b, b, size, cudaMemcpyHostToDevice );

  add<<< std::ceil(N / (double)THREADS_PER_BLOCK), THREADS_PER_BLOCK>>>( d_a, d_b, d_c );
  cudaDeviceSynchronize();


  cudaMemcpy( c, d_c, size, cudaMemcpyDeviceToHost);

  bool success = true;
  for( int i = 0; i < N; i++ )
  {
    // if( c[i] != a[i] + b[i])
    if( c[i] != a[i] )
    {
      printf("c[%d] = %d\n",i,c[i] );
      success = false;
      break;
    }
  }

  printf("%s\n", success ? "success" : "fail");

  free(a);
  free(b);
  free(c);
  cudaFree( d_a );
  cudaFree( d_b );
  cudaFree( d_c );

  return 0;
}

对于核函数：

__global__ void add(std::complex<double> *a, std::complex<double> *b, std::complex<double> *c)
    {
      int index = threadIdx.x + blockIdx.x * blockDim.x;
      // c[index] = a[index] + b[index];
      // c[index] = a[index].real();
      c[index] = a[index];
    }

线

c[index] = a[index];

会调用std::complex operator =，这可以通过编译， 但是当更改为使用行编译时：

c[index] = a[index] + b[index]; // first one
c[index] = a[index].real();     // second one

它将无法编译，第一个错误消息是：

complex.cu(10): 错误：调用 host 函数("std::operator global 函数中的 + ") ("add") 是不允许的

complex.cu(10)：错误：标识符“std::operator +”是 设备代码中未定义

改用第二个时的错误信息如下：

complex.cu(11)：错误：调用 constexpr host 函数（“real”） 不允许来自 global 函数（“add”）。实验 可以使用标志 '--expt-relaxed-constexpr' 来允许这样做。

在编译过程中检测到 1 个错误 “/tmp/tmpxft_000157af_00000000-8_complex.cpp1.ii”。

我使用的编译命令：

/usr/local/cuda-10.2/bin/nvcc -o complex complex.cu

我很清楚设备代码不能调用主机代码，而 std::complex 的 real() 和 + 函数都是主机代码，所以它们不能在内核函数中调用，但是我不明白为什么 std::复杂运算符 = 可以在我的内核函数中传递编译吗？

更新： 重载 std::complex 的 operator+ 后，上面的代码可以达到预期的效果：

__host__ __device__ std::complex<double> operator+(const std::complex<double>& a, const std::complex<double>& b)

{

    const double* aArg = reinterpret_cast<const double*>(&a);

    const double* bArg = reinterpret_cast<const double*>(&b);

    double retVal[2] = { aArg[0] + bArg[0], aArg[1] + bArg[1] };

    return std::move(*reinterpret_cast<std::complex<double>*>(retVal));
 
}

根本原因是 std::complex 的下划线结构实际上是您定义的 2 种数据类型的数组，例如 double[2]，好处是我们可以在主机/设备端拥有相同的函数参数。但是，我仍然建议在 CUDA 中使用thrust/complex 或其他复杂的库。

Answer 1

不，CUDA C++ 没有将std::complex<T>::operator+() 作为内置实现。

std::complex<T> 类型未针对 GPU 实现；它的所有方法都是仅主机编写的。例外是constexpr 方法，正如@RobertCrovella 指出的那样，编译器愿意将一些/所有隐式声明的方法视为__host__ __device__ - 例如复制构造函数或赋值运算符。这就是 c[index] = a[index] 起作用的原因：它使用隐式定义的赋值运算符。

对于在设备端使用复数，请考虑以下问题：

CUDA - How to work with complex numbers?