CUDA设备间的数据传输代价高昂

Question

我已经编写了一些代码，试图交换存储在平面数组中的二维矩阵的四个象限，以用于FFT目的。

    int leftover = W-dcW;

    T *temp;
    T *topHalf;
cudaMalloc((void **)&temp, dcW * sizeof(T));

    //swap every row, left and right
    for(int i = 0; i < H; i++)
    {
        cudaMemcpy(temp, &data[i*W], dcW*sizeof(T),cudaMemcpyDeviceToDevice);
        cudaMemcpy(&data[i*W],&data[i*W+dcW], leftover*sizeof(T), cudaMemcpyDeviceToDevice);
        cudaMemcpy(&data[i*W+leftover], temp, dcW*sizeof(T), cudaMemcpyDeviceToDevice); 
    }

cudaMalloc((void **)&topHalf, dcH*W* sizeof(T));
    leftover = H-dcH;
    cudaMemcpy(topHalf, data, dcH*W*sizeof(T), cudaMemcpyDeviceToDevice);
    cudaMemcpy(data, &data[dcH*W], leftover*W*sizeof(T), cudaMemcpyDeviceToDevice);
    cudaMemcpy(&data[leftover*W], topHalf, dcH*W*sizeof(T), cudaMemcpyDeviceToDevice);

请注意，此代码使用设备指针，并进行DeviceToDevice传输。

为什么这似乎运行得如此缓慢？是否可以进行优化？我将其与在主机上使用常规memcpy执行相同操作的时间进行了计时，它大约慢了2倍。

有什么想法吗？

- Derek

---

启动cudaMemcpy操作代价高昂。你最好编写一个从输入读取数据，交换数据并写入到适当位置的核函数，而不是把cudaMemcpy放在for循环中。 - Pavan Yalamanchili

嗯...糟糕。关于执行主机memcpy和传输到设备的比较怎么样？ - Derek

Answer 1

我最终编写了一个内核来进行交换。这比设备到设备的memcpy操作要快。

- Derek

Answer 2

也许以下在CUDA中执行2D fftshift的解决方案会引起您的兴趣：

#define IDX2R(i,j,N) (((i)*(N))+(j))

__global__ void fftshift_2D(double2 *data, int N1, int N2)
{
    int i = threadIdx.y + blockDim.y * blockIdx.y;
    int j = threadIdx.x + blockDim.x * blockIdx.x;

    if (i < N1 && j < N2) {
        double a = pow(-1.0, (i+j)&1);

        data[IDX2R(i,j,N2)].x *= a;
        data[IDX2R(i,j,N2)].y *= a;
    }
}

这里的方法是将待转换矩阵乘以一个棋盘，其中棋盘由1和-1组成，这相当于乘以exp(-j*(n+m)*pi)，从而在共轭域中实现双向移位。在应用CUFFT之前和之后必须调用此核函数。它的一个优点是避免了内存移动/交换。 速度改进 根据在 NVIDIA论坛收到的建议，通过更改指令可以提高速度。

double a = pow(-1.0,(i+j)&1);

to

double a = 1-2*((i+j)&1);

为避免使用缓慢的pow例程。

- Vitality

---

实际上，在几乎所有的过滤应用中，通过将所有的过滤器保留在包装的fft空间中，可以省略此步骤。 - Mikhail