CUDA設備內存拷貝：cudaMemcpyDeviceToDevice vs拷貝內核

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

對於普通的線性到線性內存複製，共享內存不會給您帶來任何好處。你的內核應該沒問題。在使用較少數量的線程塊運行時，可能會進行一些小的優化，但在某種程度上調整它將取決於您的特定GPU。

共享內存可用於在進行某種修改複製（如轉置操作）的內核中發揮良好效果。在這些情況下，通過共享內存訪問的成本被改進的合併性能所抵消。但是對於你的內核來說，讀寫操作都應該結合在一起。

對於單個大型複製操作，cudaMemcpyDeviceToDevice應該提供非常好的性能，因爲單個調用的開銷在整個數據移動中分攤。也許你應該用兩種方法 - 用nvprof很容易。評論中引用的討論涉及矩陣象限被交換的特定用例。在這種情況下，一個NxN矩陣需要〜1.5N cudaMemcpy操作，但與單個內核調用進行比較。在這種情況下，API調用設置的開銷將開始成爲一個重要因素。但是，將單個cudaMemcpy操作與單個等效內核調用進行比較時，cudaMemcpy操作應該很快。

__constant__內存不能被設備代碼修改，所以您將不得不使用基於cudaMemcpyFromSymbol和cudaMemcpyToSymbol的主機代碼。

Answer 2

Robert Crovella已經回答了這個問題。我在這裏只是提供了一個示例代碼爲內存拷貝兩種方法從設備到設備的CUDA比較：

1. 使用cudaMemcpyDeviceToDevice;
2. 使用複製內核。

THE CODE

測試代碼如下：

#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文件保持here，而 的TimingGPU.cu和TimingGPU.cuh保持here。一塊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 

結果證實羅伯特Crovella猜想：cudaMemcpyDeviceToDevice一般都在拷貝內核較好。