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

Question

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

这是一个naive kernel:

__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上）大于常量内存，如何将其复制到具有__const__的另一个GPU内存中？

- Zziggurats

---

你认为为什么cudamemcpyDeviceToDevice的性能较差？ - talonmies

http://stackoverflow.com/questions/22284533/cuda-least-square-solving-poor-in-speed 中的评论 - Zziggurats

我认为你不需要共享内存。你必须逐个元素地复制两个数组，并且在第一个核函数中没有线程协作。 - Vitality

第一个内核就足够了吗？ - Zziggurats

Answer 1

Robert Crovella已经回答了这个问题。我在这里提供一个示例代码，用于比较在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 通常比使用拷贝内核更好。

- Vitality

Answer 2

对于普通的线性到线性的内存复制，共享内存不会给您带来任何好处。你的简单kernel应该没问题。在运行较少的线程块方面可能有一些小的优化措施，但这将在一定程度上取决于您特定的GPU。

共享内存可以在进行某种修改复制（例如转置操作）的kernel中得到很好的效果。在这些情况下，通过共享内存的传输成本被改善的合并执行性能所抵消。但是对于您的naive kernel，读和写都应该合并。

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

__constant__内存不能被设备代码修改，因此您将不得不使用基于cudaMemcpyFromSymbol和cudaMemcpyToSymbol的host代码。

- Robert Crovella

Answer 3

#include <iostream>
#include <vector>
#include <iomanip>
#include <cuda_runtime.h>

#define CHECK_CUDA(cond) check_cuda(cond, __LINE__)

void check_cuda(cudaError_t status, std::size_t line)
{
    if(status != cudaSuccess)
    {
        std::cout << cudaGetErrorString(status) << '\n';
        std::cout << "Line: " << line << '\n';
        throw 0;
    }
}

__global__ void copy_kernel(float* __restrict__ output, const float* __restrict__ input, int N)
{
    for (int i = blockIdx.x * blockDim.x + threadIdx.x;  i < N; i += blockDim.x * gridDim.x) 
        output[i] = input[i];
}

int main()
{
    constexpr int num_trials = 100;
    std::vector<int> test_sizes = { 100'000, 1'000'000, 10'000'000, 100'000'000, 250'000'000 };

    int grid_size = 0, block_size = 0;
    CHECK_CUDA(cudaOccupancyMaxPotentialBlockSize(&grid_size, &block_size, copy_kernel, 0));

    std::cout << std::fixed << std::setprecision(4) << std::endl;

    for (auto sz : test_sizes)
    {
        std::cout << "Test Size: " << sz << '\n';

        float *d_vector_src = nullptr, *d_vector_dest = nullptr;
        CHECK_CUDA(cudaMalloc(&d_vector_src, sz * sizeof(float)));
        CHECK_CUDA(cudaMalloc(&d_vector_dest, sz * sizeof(float)));

        cudaEvent_t start, stop;
        CHECK_CUDA(cudaEventCreate(&start));
        CHECK_CUDA(cudaEventCreate(&stop));

        float accumulate = 0.0;
        for (int i = 0; i < num_trials; i++)
        {
            CHECK_CUDA(cudaEventRecord(start));
            copy_kernel<<<grid_size, block_size>>>(d_vector_dest, d_vector_src, sz);
            CHECK_CUDA(cudaEventRecord(stop));
            CHECK_CUDA(cudaEventSynchronize(stop));

            float current_time = 0;
            CHECK_CUDA(cudaEventElapsedTime(&current_time, start, stop));
            accumulate += current_time;
        }
        std::cout << "\tKernel Copy Time: " << accumulate / num_trials << "ms\n";

        accumulate = 0.0;
        for (int i = 0; i < num_trials; i++)
        {
            CHECK_CUDA(cudaEventRecord(start));
            CHECK_CUDA(cudaMemcpy(d_vector_dest, d_vector_src, sz * sizeof(float), cudaMemcpyDeviceToDevice));
            CHECK_CUDA(cudaEventRecord(stop));
            CHECK_CUDA(cudaEventSynchronize(stop));

            float current_time = 0;
            CHECK_CUDA(cudaEventElapsedTime(&current_time, start, stop));
            accumulate += current_time;
        }
        std::cout << "\tMemcpy Time: " << accumulate / num_trials << "ms\n";

        CHECK_CUDA(cudaFree(d_vector_src));
        CHECK_CUDA(cudaFree(d_vector_dest));
    }

    return 0;
}

移动版GTX 1050

Test Size: 100000
        Kernel Copy Time: 0.0118ms
        Memcpy Time: 0.0127ms
Test Size: 1000000
        Kernel Copy Time: 0.0891ms
        Memcpy Time: 0.0899ms
Test Size: 10000000
        Kernel Copy Time: 0.8697ms
        Memcpy Time: 0.8261ms
Test Size: 100000000
        Kernel Copy Time: 8.8871ms
        Memcpy Time: 8.2401ms
Test Size: 250000000
        Kernel Copy Time: 22.3060ms
        Memcpy Time: 20.5419ms

GTX 1080 Ti

Test Size: 100000
    Kernel Copy Time: 0.0166ms
    Memcpy Time: 0.0188ms
Test Size: 1000000
    Kernel Copy Time: 0.0580ms
    Memcpy Time: 0.0727ms
Test Size: 10000000
    Kernel Copy Time: 0.4674ms
    Memcpy Time: 0.5047ms
Test Size: 100000000
    Kernel Copy Time: 4.7992ms
    Memcpy Time: 3.7722ms
Test Size: 250000000
    Kernel Copy Time: 7.2485ms
    Memcpy Time: 5.5863ms
Test Size: 1000000000
    Kernel Copy Time: 31.5570ms
    Memcpy Time: 22.3184ms

RTX 2080 Ti

Test Size: 100000
    Kernel Copy Time: 0.0048ms
    Memcpy Time: 0.0054ms
Test Size: 1000000
    Kernel Copy Time: 0.0193ms
    Memcpy Time: 0.0220ms
Test Size: 10000000
    Kernel Copy Time: 0.1578ms
    Memcpy Time: 0.1537ms
Test Size: 100000000
    Kernel Copy Time: 2.1156ms
    Memcpy Time: 1.5006ms
Test Size: 250000000
    Kernel Copy Time: 5.5195ms
    Memcpy Time: 3.7424ms
Test Size: 1000000000
    Kernel Copy Time: 23.2106ms
    Memcpy Time: 14.9483ms

- Yashas

---

这是一个很棒的回答。我点了赞。我看了你的网络档案，你在物理和化学方面有很多积分。我想知道你是否支持我的提议，建立一个材料建模Stack Exchange。我一直在努力，但我们仍需要更多的承诺者。 - Nike