CUDA 设备内存副本：cudaMemcpyDeviceToDevice 与复制内核

Question

我正在编写一个 cuda 内核来将一个数组复制到另一个数组。它们都在 GPU 内存中。我不想使用，cudamemcpyDeviceToDevice因为它的性能很差。

天真的内核：

__global__ void GpuCopy( float* des , float* __restrict__ sour ,const int M , const int N )
{
    int tx=blockIdx.x*blockDim.x+threadIdx.x;
    if(tx<N*M)
        des[tx]=sour[tx];
}   

我认为朴素的内核不会获得高性能，所以我尝试使用__shared__内存，但看起来不太好：

__shared__ float TILE[tile];
int tid=threadIdx.x;
for(int i=0; i<M*N/tile;i++)
{
    TILE[tid]=sour[i*tile+tid]
    des[i*tile+tid]=TILE[tid]
}

前一个代码片段将全局内存复制到des[]，而后者将全局内存复制到__shared__然后再复制__shared__到des[]。我认为后者比前者慢。

那么，如何编写__shared__复制内存的代码呢？另一个问题是，如果我想使用__const__内存并且数组（已经在 GPU 中）大于常量内存，如何将其复制到另一个 GPU 内存__const__？

Answer 1

罗伯特·克罗维拉已经回答了这个问题。我在这里只是提供一个示例代码来比较 CUDA 中从设备到设备的内存复制的两种方法：

1. 使用cudaMemcpyDeviceToDevice;
2. 使用复制内核。

代码

测试代码如下：

#include <stdio.h>

#include "Utilities.cuh"
#include "TimingGPU.cuh"

#define BLOCKSIZE   512

/***************/
/* COPY KERNEL */
/***************/
__global__ void copyKernel(const double * __restrict__ d_in, double * __restrict__ d_out, const int N) {

    const int tid = threadIdx.x + blockIdx.x * blockDim.x;

    if (tid >= N) return;

    d_out[tid] = d_in[tid];

}

/********/
/* MAIN */
/********/
int main() {

    const int N = 1000000;

    TimingGPU timerGPU;

    double *h_test = (double *)malloc(N * sizeof(double));

    for (int k = 0; k < N; k++) h_test[k] = 1.;

    double *d_in;   gpuErrchk(cudaMalloc(&d_in, N * sizeof(double)));
    gpuErrchk(cudaMemcpy(d_in, h_test, N * sizeof(double), cudaMemcpyHostToDevice));

    double *d_out; gpuErrchk(cudaMalloc(&d_out, N * sizeof(double)));

    timerGPU.StartCounter();
    gpuErrchk(cudaMemcpy(d_out, d_in, N * sizeof(double), cudaMemcpyDeviceToDevice));
    printf("cudaMemcpy timing = %f [ms]\n", timerGPU.GetCounter());

    timerGPU.StartCounter();
    copyKernel << <iDivUp(N, BLOCKSIZE), BLOCKSIZE >> >(d_in, d_out, N);
    gpuErrchk(cudaPeekAtLastError());
    gpuErrchk(cudaDeviceSynchronize());
    printf("Copy kernel timing = %f [ms]\n", timerGPU.GetCounter());

    return 0;
}

和文件Utilities.cu都Utilities.cuh保存在这里，而 和TimingGPU.cu则TimingGPU.cuh保存在这里。

时机

在 GeForce GTX960 卡上进行的测试。时间以毫秒为单位。

N           cudaMemcpyDeviceToDevice           copy kernel
1000        0.0075                             0.029
10000       0.0078                             0.072
100000      0.019                              0.068
1000000     0.20                               0.22

结果证实了 Robert Crovella 的猜想：cudaMemcpyDeviceToDevice通常比复制内核更可取。

Answer 2

For ordinary linear-to-linear memory copying, shared memory won't give you any benefit. Your naive kernel should be fine. There may be some small optimizations that could be made in terms of running with a smaller number of threadblocks, but tuning this will be dependent on your specific GPU, to some degree.

Shared memory can be used to good effect in kernels that do some kind of modified copying, such as a transpose operation. In these cases, the cost of the trip through shared memory is offset by the improved coalescing performance. But with your naive kernel, both reads and writes should coalesce.

对于单个大型复制操作，cudaMemcpyDeviceToDevice应该提供非常好的性能，因为单个调用的开销在整个数据移动中分摊。也许您应该对这两种方法进行计时——这很容易做到nvprof。评论中引用的讨论涉及交换矩阵象限的特定用例。在这种情况下，NxN 矩阵需要约 1.5N 次cudaMemcpy操作，但与单个内核调用进行比较。在这种情况下，API 调用设置的开销将开始成为一个重要因素。但是，当将单个cudaMemcpy操作与单个等效内核调用进行比较时，该cudaMemcpy操作应该很快。

__constant__设备代码无法修改内存，因此您必须使用基于cudaMemcpyFromSymbol和 的主机代码cudaMemcpyToSymbol。