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

Question

我已经将向量add的int版本修改为两个复数向量相加，下面的代码可以工作，但我很困惑：

#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运算符=，这样可以通过编译， 但是当更改为使用以下行进行编译时：

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::complex 运算符= 可以在我的内核函数中通过编译吗？

更新： 重载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 - 如何处理复数？