在cuda上加載矢量的步驟

Question

我想在CUDA C++代碼上運行矢量加法函數，但對於大小爲5,000,000的大型浮點數組，它的運行速度比我的CPU版本慢。下面是我說的相關CUDA和CPU代碼：在比vectorStepAdd3計算慢

#define THREADS_PER_BLOCK 1024 
typedef float real; 
__global__ void vectorStepAddKernel2(real*x, real*y, real*z, real alpha, real beta, int size, int xstep, int ystep, int zstep) 
{ 
    int i = blockDim.x * blockIdx.x + threadIdx.x; 
    if (i < size) 
    { 
     x[i*xstep] = alpha* y[i*ystep] + beta*z[i*zstep]; 
    } 
} 

cudaError_t vectorStepAdd2(real *x, real*y, real* z, real alpha, real beta, int size, int xstep, int ystep, int zstep) 
{ 

    cudaError_t cudaStatus; 
    int threadsPerBlock = THREADS_PER_BLOCK; 
    int blocksPerGrid = (size + threadsPerBlock -1)/threadsPerBlock; 
    vectorStepAddKernel2<<<blocksPerGrid, threadsPerBlock>>>(x, y, z, alpha, beta, size, xstep, ystep, zstep); 

    // cudaDeviceSynchronize waits for the kernel to finish, and returns 
    // any errors encountered during the launch. 
    cudaStatus = cudaDeviceSynchronize(); 
    if (cudaStatus != cudaSuccess) { 
     fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching vectorStepAddKernel!\n", cudaStatus); 
     exit(1); 
    } 

    return cudaStatus; 
} 

//CPU function: 

void vectorStepAdd3(real *x, real*y, real* z, real alpha, real beta, int size, int xstep, int ystep, int zstep) 
{ 
    for(int i=0;i<size;i++) 
    { 
     x[i*xstep] = alpha* y[i*ystep] + beta*z[i*zstep]; 
    } 
} 

調用vectorStepAdd2結果時各3個數組的是大小5000000和大小= 50000（即，50000個元素添加一起以這種逐步的方式）。

關於我能做些什麼來加速GPU代碼的任何想法？ 我的設備是特斯拉M2090 GPU

感謝

Answer 1

在回答你的問題「什麼我可以做，以加快GPU代碼的任何想法？」

首先，讓我先說一下這樣的陳述，即所提出的操作X = alpha * Y + beta * Z沒有每個字節所需的數據傳輸的大量計算強度。因此，我無法在這個特定的代碼上擊敗CPU時間。然而，它可能是有益的覆蓋2個思路，加快驗證碼：

1. 使用page-locked內存進行數據傳輸操作。這爲GPU版本的數據傳輸時間減少了約2倍，GPU版本佔據了整個GPU執行時間的主導地位。

2. 按照@njuffa here的建議，使用cudaMemcpy2D的跨步複製技術。結果是2倍：我們可以將數據傳輸量減少到計算所需的數據量，即和，然後我們可以重新編寫內核來對註釋中建議的數據進行連續操作（同樣由njuffa）。這使得數據傳輸時間進一步提高了3倍，內核計算時間提高了10倍。

此代碼提供了這些操作的示例：

#include <stdio.h> 
#include <stdlib.h> 


#define THREADS_PER_BLOCK 1024 
#define DSIZE 5000000 
#define WSIZE 50000 
#define XSTEP 47 
#define YSTEP 43 
#define ZSTEP 41 
#define TOL 0.00001f 


#define cudaCheckErrors(msg) \ 
    do { \ 
     cudaError_t __err = cudaGetLastError(); \ 
     if (__err != cudaSuccess) { \ 
      fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \ 
       msg, cudaGetErrorString(__err), \ 
       __FILE__, __LINE__); \ 
      fprintf(stderr, "*** FAILED - ABORTING\n"); \ 
      exit(1); \ 
     } \ 
    } while (0) 

typedef float real; 

__global__ void vectorStepAddKernel2(real *x, real *y, real *z, real alpha, real beta, int size, int xstep, int ystep, int zstep) 
{ 
    int i = blockDim.x * blockIdx.x + threadIdx.x; 
    if (i < size) 
    { 
     x[i*xstep] = alpha* y[i*ystep] + beta*z[i*zstep]; 
    } 
} 

__global__ void vectorStepAddKernel2i(real *x, real *y, real *z, real alpha, real beta, int size) 
{ 
    int i = blockDim.x * blockIdx.x + threadIdx.x; 
    if (i < size) 
    { 
     x[i] = alpha* y[i] + beta*z[i]; 
    } 
} 

void vectorStepAdd2(real *x, real *y, real *z, real alpha, real beta, int size, int xstep, int ystep, int zstep) 
{ 

    int threadsPerBlock = THREADS_PER_BLOCK; 
    int blocksPerGrid = (size + threadsPerBlock -1)/threadsPerBlock; 
    vectorStepAddKernel2<<<blocksPerGrid, threadsPerBlock>>>(x, y, z, alpha, beta, size, xstep, ystep, zstep); 
    cudaDeviceSynchronize(); 
    cudaCheckErrors("kernel2 fail"); 
} 


void vectorStepAdd2i(real *x, real *y, real *z, real alpha, real beta, int size) 
{ 

    int threadsPerBlock = THREADS_PER_BLOCK; 
    int blocksPerGrid = (size + threadsPerBlock -1)/threadsPerBlock; 
    vectorStepAddKernel2i<<<blocksPerGrid, threadsPerBlock>>>(x, y, z, alpha, beta, size); 
    cudaDeviceSynchronize(); 
    cudaCheckErrors("kernel3 fail"); 
} 

//CPU function: 

void vectorStepAdd3(real *x, real*y, real* z, real alpha, real beta, int size, int xstep, int ystep, int zstep) 
{ 
    for(int i=0;i<size;i++) 
    { 
     x[i*xstep] = alpha* y[i*ystep] + beta*z[i*zstep]; 
    } 
} 

int main() { 

    real *h_x, *h_y, *h_z, *c_x, *h_x1; 
    real *d_x, *d_y, *d_z, *d_x1, *d_y1, *d_z1; 

    int dsize = DSIZE; 
    int wsize = WSIZE; 
    int xstep = XSTEP; 
    int ystep = YSTEP; 
    int zstep = ZSTEP; 
    real alpha = 0.5f; 
    real beta = 0.5f; 
    float et; 

/* 
    h_x = (real *)malloc(dsize*sizeof(real)); 
    if (h_x == 0){printf("malloc1 fail\n"); return 1;} 
    h_y = (real *)malloc(dsize*sizeof(real)); 
    if (h_y == 0){printf("malloc2 fail\n"); return 1;} 
    h_z = (real *)malloc(dsize*sizeof(real)); 
    if (h_z == 0){printf("malloc3 fail\n"); return 1;} 
    c_x = (real *)malloc(dsize*sizeof(real)); 
    if (c_x == 0){printf("malloc4 fail\n"); return 1;} 
    h_x1 = (real *)malloc(dsize*sizeof(real)); 
    if (h_x1 == 0){printf("malloc1 fail\n"); return 1;} 
*/ 

    cudaHostAlloc((void **)&h_x, dsize*sizeof(real), cudaHostAllocDefault); 
    cudaCheckErrors("cuda Host Alloc 1 fail"); 
    cudaHostAlloc((void **)&h_y, dsize*sizeof(real), cudaHostAllocDefault); 
    cudaCheckErrors("cuda Host Alloc 2 fail"); 
    cudaHostAlloc((void **)&h_z, dsize*sizeof(real), cudaHostAllocDefault); 
    cudaCheckErrors("cuda Host Alloc 3 fail"); 
    cudaHostAlloc((void **)&c_x, dsize*sizeof(real), cudaHostAllocDefault); 
    cudaCheckErrors("cuda Host Alloc 4 fail"); 
    cudaHostAlloc((void **)&h_x1, dsize*sizeof(real), cudaHostAllocDefault); 
    cudaCheckErrors("cuda Host Alloc 5 fail"); 


    cudaMalloc((void **)&d_x, dsize*sizeof(real)); 
    cudaCheckErrors("cuda malloc1 fail"); 
    cudaMalloc((void **)&d_y, dsize*sizeof(real)); 
    cudaCheckErrors("cuda malloc2 fail"); 
    cudaMalloc((void **)&d_z, dsize*sizeof(real)); 
    cudaCheckErrors("cuda malloc3 fail"); 
    cudaMalloc((void **)&d_x1, wsize*sizeof(real)); 
    cudaCheckErrors("cuda malloc4 fail"); 
    cudaMalloc((void **)&d_y1, wsize*sizeof(real)); 
    cudaCheckErrors("cuda malloc5 fail"); 
    cudaMalloc((void **)&d_z1, wsize*sizeof(real)); 
    cudaCheckErrors("cuda malloc6 fail"); 

    for (int i=0; i< dsize; i++){ 
    h_x[i] = 0.0f; 
    h_x1[i] = 0.0f; 
    c_x[i] = 0.0f; 
    h_y[i] = (real)(rand()/(real)RAND_MAX); 
    h_z[i] = (real)(rand()/(real)RAND_MAX); 
    } 


    cudaEvent_t t_start, t_stop, k_start, k_stop; 
    cudaEventCreate(&t_start); 
    cudaEventCreate(&t_stop); 
    cudaEventCreate(&k_start); 
    cudaEventCreate(&k_stop); 
    cudaCheckErrors("event fail"); 

    // first test original GPU version 

    cudaEventRecord(t_start); 
    cudaMemcpy(d_x, h_x, dsize * sizeof(real), cudaMemcpyHostToDevice); 
    cudaCheckErrors("cuda memcpy 1 fail"); 
    cudaMemcpy(d_y, h_y, dsize * sizeof(real), cudaMemcpyHostToDevice); 
    cudaCheckErrors("cuda memcpy 2 fail"); 
    cudaMemcpy(d_z, h_z, dsize * sizeof(real), cudaMemcpyHostToDevice); 
    cudaCheckErrors("cuda memcpy 3 fail"); 


    cudaEventRecord(k_start); 
    vectorStepAdd2(d_x, d_y, d_z, alpha, beta, wsize, xstep, ystep, zstep); 
    cudaEventRecord(k_stop); 

    cudaMemcpy(h_x, d_x, dsize * sizeof(real), cudaMemcpyDeviceToHost); 
    cudaCheckErrors("cuda memcpy 4 fail"); 
    cudaEventRecord(t_stop); 
    cudaEventSynchronize(t_stop); 
    cudaEventElapsedTime(&et, t_start, t_stop); 
    printf("GPU original version total elapsed time is: %f ms.\n", et); 
    cudaEventElapsedTime(&et, k_start, k_stop); 
    printf("GPU original kernel elapsed time is: %f ms.\n", et); 

    //now test CPU version 

    cudaEventRecord(t_start); 
    vectorStepAdd3(c_x, h_y, h_z, alpha, beta, wsize, xstep, ystep, zstep); 
    cudaEventRecord(t_stop); 
    cudaEventSynchronize(t_stop); 
    cudaEventElapsedTime(&et, t_start, t_stop); 
    printf("CPU version total elapsed time is: %f ms.\n", et); 
    for (int i = 0; i< dsize; i++) 
    if (fabsf((float)(h_x[i]-c_x[i])) > TOL) { 
     printf("cpu/gpu results mismatch at i = %d, cpu = %f, gpu = %f\n", i, c_x[i], h_x[i]); 
     return 1; 
     } 


    // now test improved GPU version 

    cudaEventRecord(t_start); 
// cudaMemcpy2D(d_x1, sizeof(real), h_x, xstep * sizeof(real), sizeof(real), wsize, cudaMemcpyHostToDevice); 
// cudaCheckErrors("cuda memcpy 5 fail"); 
    cudaMemcpy2D(d_y1, sizeof(real), h_y, ystep * sizeof(real), sizeof(real), wsize, cudaMemcpyHostToDevice); 
    cudaCheckErrors("cuda memcpy 6 fail"); 
    cudaMemcpy2D(d_z1, sizeof(real), h_z, zstep * sizeof(real), sizeof(real), wsize, cudaMemcpyHostToDevice); 
    cudaCheckErrors("cuda memcpy 7 fail"); 

    cudaEventRecord(k_start); 
    vectorStepAdd2i(d_x1, d_y1, d_z1, alpha, beta, wsize); 
    cudaEventRecord(k_stop); 

    cudaMemcpy2D(h_x1, xstep*sizeof(real), d_x1, sizeof(real), sizeof(real), wsize, cudaMemcpyDeviceToHost); 
    cudaCheckErrors("cuda memcpy 8 fail"); 
    cudaEventRecord(t_stop); 
    cudaEventSynchronize(t_stop); 
    cudaEventElapsedTime(&et, t_start, t_stop); 
    printf("GPU improved version total elapsed time is: %f ms.\n", et); 
    cudaEventElapsedTime(&et, k_start, k_stop); 
    printf("GPU improved kernel elapsed time is: %f ms.\n", et); 

    for (int i = 0; i< dsize; i++) 
    if (fabsf((float)(h_x[i]-h_x1[i])) > TOL) { 
     printf("gpu/gpu improved results mismatch at i = %d, gpu = %f, gpu imp = %f\n", i, h_x[i], h_x1[i]); 
     return 1; 
     } 

    printf("Results:i CPU  GPU  GPUi \n"); 
    for (int i = 0; i< 20*xstep; i+=xstep) 
    printf(" %d   %f  %f  %f %f %f\n",i, c_x[i], h_x[i], h_x1[i]); 


    return 0; 
} 

如前所述，我還是沒能擊敗的CPU時間，而我認爲這要麼我自己缺乏編碼技能或否則這個操作基本上沒有足夠的計算複雜度在GPU上感興趣。不過這裏有一些樣品的結果：

GPU original version total elapsed time is: 13.352256 ms. 
GPU original kernel elapsed time is: 0.195808 ms. 
CPU version total elapsed time is: 2.599584 ms. 
GPU improved version total elapsed time is: 4.228288 ms. 
GPU improved kernel elapsed time is: 0.027392 ms. 
Results:i CPU  GPU  GPUi 
    0   0.617285  0.617285  0.617285 
    47   0.554522  0.554522  0.554522 
    94   0.104245  0.104245  0.104245 
.... 

我們可以看到，改進後的內核具有約3倍的整體減少相比原來的內核，幾乎全部是由於數據的減少拷貝時間。數據複製時間的縮短是由於改進後的2D memcpy，我們只需要複製實際使用的數據。 （沒有頁面鎖定的內存，這些數據傳輸時間大約是其兩倍）。我們還可以看到內核計算時間比原始內核的CPU計算速度快10倍，比改進內核的CPU計算快100倍。不過，考慮到數據傳輸時間，我們無法克服CPU的速度。

最後一點評論是，cudaMemcpy2D操作的「成本」仍然很高。爲了減少100倍的矢量大小，我們只能看到複製時間縮短3倍。因此，跨越式訪問仍然會導致使用GPU的相對昂貴的方式。如果我們只是傳輸50,000個連續元素的向量，我們預計複製時間幾乎會線性減少100倍（與5000000個元素的原始複製向量相比）。這意味着複製的時間將少於1毫秒，我們的GPU版本將比CPU更快，至少這個天真的單線程CPU代碼。