淺談CUDA 平行運算機制之 cudaMemcpy

完成【AI人工智慧/機器、深度學習/平行運算:模組A】 CUDA 平行運算技術與實作實戰班課程訓練，我在課堂上加強強調GPU的運算能力是超級強大，但CUDA是異質多核心系統架構，在CPU是x86架構，GPU是NVIDIA的，採PCIe BUS的架構下，因為PCIe BUS架構 的物理特性限制，一個標準的PCI-E 1.0 X16，匯流排寬度16位元，工作時脈2.5 GHz，其資料速率是8 GB/s（雙工），但實際上通常為5-6GB/s。
一般大多數採CUDA平行運算應用程式，是一開始做一次cudaMalloc，之後不再調用它了。
之後會在x86 Host端跟GPU Kernel端透過cudaMemcpy 在PCIe BUS間互傳資料，所以整體的CUDA平行運算應用程式其瓶頸會在cudaMemcpy函式的呼叫上。 現簡單驗證:
1. 循序執行程式如下，numElements是100000000，程式是將numElements個的float type矩陣相加：
cpu_startTime = clock();
for (int i = 0; i < numElements; ++i)
{
h_C[i] = h_A[i] + h_B[i];
}
cpu_endTime = clock();
執行所花的時間是 Sequential cpu_endTime - cpu_startTime=273.000000。

2. CUDA 平行運算程式如下：
/**
 * CUDA Kernel Device code
 * Computes the vector addition of A and B into C. The 3 vectors have the same
 * number of elements numElements.
 */
__global__ void vectorAdd(const float *A, const float *B, float *C, int numElements)
{
    int i = blockDim.x * blockIdx.x + threadIdx.x;

    if (i < numElements)
    {
        C[i] = A[i] + B[i];
    }
}
GPU 使用 390625 blocks，每個Block 配置 256 threads，測定CUDA 平行運算執行程式執行所花的時間，包含執行GPU CUDA的vectorAdd Kernel程式 與 將結果複製到 x86 Host端，整個所需的時間(PS:如果單測vectorAdd Kernel程式一定是超級快的)，程式片段如下:

// Launch the Vector Add CUDA Kernel
cpu_startTime1 = clock();
vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, numElements);
err = cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
cpu_endTime1 = clock();

執行所花的時間是 Parallel cpu_endTime1 - cpu_startTime1=167.000000。
cudaMemcpy要複製的memory size是100000000個的float變數，大約380MByte。
若PCI-E 1.0 X16實際資料速率是5 GB/s，那 1ms大約是5M Byte，一對比下來，瓶頸真的是cudaMemcpy了。
執行畫面如下: