Шаг векторного добавления медленнее на cuda

Question

Я пытаюсь запустить функцию добавления дополнительных шагов на CUDA C++, но для больших массивов с плавающей запятой размером 5 000 000 он работает медленнее, чем моя версия процессора. Ниже приводится соответствующая CUDA и процессор код, который я имею в виду:

#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]; 
    } 
} 

Calling результатов vectorStepAdd2 более медленных вычислений, чем vectorStepAdd3, когда каждые из 3-х массивов размера 5000000 и размера = 50000 (т.е. добавляются 50000 элементов вместе это поэтапно).

Любые идеи о том, что я могу сделать, чтобы ускорить код графического процессора? Мое устройство является Tesla M2090 GPU

Благодаря

Answer 1

В ответ на ваш вопрос «Любые идеи о том, что я могу сделать, чтобы ускорить код GPU?»

Прежде всего позвольте мне изложить это с утверждением, что предлагаемая операция X = alpha * Y + beta * Z не имеет большого количества интенсивности вычислений на каждый бит передачи данных. В результате я не смог победить процессорное время на этом конкретном коде. Однако это может оказаться полезным, чтобы покрыть 2 идеи, чтобы ускорить этот код:

1. Использование page-locked памяти для операций передачи данных. Это привело к сокращению примерно на 2 раза за время передачи данных для версии GPU, которая доминировала в общем времени исполнения для версии GPU.

2. Использовать технологию скользящего копирования с помощью cudaMemcpy2D, предложенную @njuffa here. Результат в 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; 
} 

Как уже упоминалось, я до сих пор не смог обыграть процессорное время, и я отношу это ни к моему собственному отсутствие кодирования навыков или иначе тот факт, что эта операция в принципе не имеет достаточной вычислительной сложности, чтобы быть интересной на графическом процессоре. Тем не менее, вот некоторые результаты выборки:

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 
.... 

Мы можем видеть, что улучшенное ядро ​​было общее сокращение около 3x по сравнению с первоначальным ядром, почти все из которых было связано с сокращением данных копирования времени. Это сокращение времени копирования данных было связано с тем, что с улучшенной 2D memcpy нам нужно только скопировать данные, которые мы фактически используем. (без памяти с записями страниц это время передачи данных будет в два раза длиннее, примерно). Мы также можем видеть, что время вычисления ядра примерно в 10 раз быстрее, чем вычисление ЦП для исходного ядра, и примерно в 100 раз быстрее, чем вычисление ЦП для улучшенного ядра. Тем не менее, с учетом времени передачи данных, мы не можем преодолеть скорость процессора.

Последний комментарий: «Стоимость» операции cudaMemcpy2D по-прежнему довольно высока. Для уменьшения 100x в размере вектора мы видим только 3-кратное сокращение времени для копирования. Таким образом, чередующийся доступ все еще делает относительно дорогой способ использования графического процессора.Если бы мы просто переносили векторы из 50 000 смежных элементов, мы ожидали бы почти линейное сокращение на 100 раз во время копирования (по сравнению с исходными векторами копирования 5000000 элементов). Это означает, что время копирования будет меньше 1 мс, а наша версия GPU будет быстрее, чем процессор, по крайней мере, этот наивный однопоточный код процессора.