Как найти сумму массива в CUDA путем уменьшения

Question

Я реализую функцию, чтобы найти сумму массива с помощью сокращения, мой массив имеет 32 * 32 элемента, а его значения равны 0 ... 1023. моя ожидаемая сумма составляет 523776, но мой результат - 15872, это неправильно. Вот мой код:

#include <stdio.h> 
#include <cuda.h> 

#define w 32 
#define h 32 
#define N w*h 

__global__ void reduce(int *g_idata, int *g_odata); 
void fill_array (int *a, int n); 

int main(void) { 
    int a[N], b[N]; // copies of a, b, c 
    int *dev_a, *dev_b; // device copies of a, b, c 
    int size = N * sizeof(int); // we need space for 512 integers 

    // allocate device copies of a, b, c 
    cudaMalloc((void**)&dev_a, size); 
    cudaMalloc((void**)&dev_b, size); 

    fill_array(a, N); 

    // copy inputs to device 
    cudaMemcpy(dev_a, a, size, cudaMemcpyHostToDevice); 
    cudaMemcpy(dev_b, b, size, cudaMemcpyHostToDevice); 

    dim3 blocksize(16,16); 
    dim3 gridsize; 

    gridsize.x=(w+blocksize.x-1)/blocksize.x; 
    gridsize.y=(h+blocksize.y-1)/blocksize.y; 

    reduce<<<gridsize, blocksize>>>(dev_a, dev_b); 

    // copy device result back to host copy of c 
    cudaMemcpy(b, dev_b, sizeof(int) , cudaMemcpyDeviceToHost); 

    printf("Reduced sum of Array elements = %d \n", b[0]); 

    cudaFree(dev_a); 
    cudaFree(dev_b); 

    return 0; 
} 

__global__ void reduce(int *g_idata, int *g_odata) { 

    __shared__ int sdata[256]; 

    // each thread loads one element from global to shared mem 
    int i = blockIdx.x*blockDim.x + threadIdx.x; 

    sdata[threadIdx.x] = g_idata[i]; 

    __syncthreads(); 
    // do reduction in shared mem 
    for (int s=1; s < blockDim.x; s *=2) 
    { 
     int index = 2 * s * threadIdx.x;; 

     if (index < blockDim.x) 
     { 
      sdata[index] += sdata[index + s]; 
     } 
     __syncthreads(); 
    } 

    // write result for this block to global mem 
    if (threadIdx.x == 0) 
     atomicAdd(g_odata,sdata[0]); 
} 

// CPU function to generate a vector of random integers 
void fill_array (int *a, int n) 
{ 
    for (int i = 0; i < n; i++) 
     a[i] = i; 
} 

Answer 1

Есть по крайней мере, 2 проблемы в коде

1. Вы делаете atomicAdd к первому элементу в вашем dev_b массиве, но не инициализирует этот элемент в известное значение (т.е. 0). Конечно, перед запуском ядра вы копируете b в dev_b, но поскольку вы не инициализировали b любым известным значениям, это не поможет. Массив b не инициализируется автоматически на ноль на C или C++, если это то, о чем вы думали. Мы можем исправить это, установив b[0] в ноль, прежде чем копировать b в dev_b.

2. Ваше редукционное ядро ​​написано для обработки 1D-кода (т. Е. Используемый только индекс потока - это 1D-индекс потока, основанный на значениях .x), но вы запускаете ядро ​​с 2D-потоковыми блоками и сетками. Это несоответствие не будет работать должным образом, и нам нужно либо запустить 1D threadblock и grid, либо перезаписать ядро ​​для работы с двумерными индексами (то есть .x и .y). Я выбрал первый (1D).

Вот обработанный пример с этими изменениями в коду, это, кажется, для получения правильного результата:

$ cat t1218.cu 
#include <stdio.h> 

#define w 32 
#define h 32 
#define N w*h 

__global__ void reduce(int *g_idata, int *g_odata); 
void fill_array (int *a, int n); 

int main(void) { 
    int a[N], b[N]; // copies of a, b, c 
    int *dev_a, *dev_b; // device copies of a, b, c 
    int size = N * sizeof(int); // we need space for 512 integers 

    // allocate device copies of a, b, c 
    cudaMalloc((void**)&dev_a, size); 
    cudaMalloc((void**)&dev_b, size); 

    fill_array(a, N); 
    b[0] = 0; //initialize the first value of b to zero 
    // copy inputs to device 
    cudaMemcpy(dev_a, a, size, cudaMemcpyHostToDevice); 
    cudaMemcpy(dev_b, b, size, cudaMemcpyHostToDevice); 

    dim3 blocksize(256); // create 1D threadblock 
    dim3 gridsize(N/blocksize.x); //create 1D grid 

    reduce<<<gridsize, blocksize>>>(dev_a, dev_b); 

    // copy device result back to host copy of c 
    cudaMemcpy(b, dev_b, sizeof(int) , cudaMemcpyDeviceToHost); 

    printf("Reduced sum of Array elements = %d \n", b[0]); 
    printf("Value should be: %d \n", ((N-1)*(N/2))); 
    cudaFree(dev_a); 
    cudaFree(dev_b); 

    return 0; 
} 

__global__ void reduce(int *g_idata, int *g_odata) { 

    __shared__ int sdata[256]; 

    // each thread loads one element from global to shared mem 
    // note use of 1D thread indices (only) in this kernel 
    int i = blockIdx.x*blockDim.x + threadIdx.x; 

    sdata[threadIdx.x] = g_idata[i]; 

    __syncthreads(); 
    // do reduction in shared mem 
    for (int s=1; s < blockDim.x; s *=2) 
    { 
     int index = 2 * s * threadIdx.x;; 

     if (index < blockDim.x) 
     { 
      sdata[index] += sdata[index + s]; 
     } 
     __syncthreads(); 
    } 

    // write result for this block to global mem 
    if (threadIdx.x == 0) 
     atomicAdd(g_odata,sdata[0]); 
} 

// CPU function to generate a vector of random integers 
void fill_array (int *a, int n) 
{ 
    for (int i = 0; i < n; i++) 
     a[i] = i; 
} 
$ nvcc -o t1218 t1218.cu 
$ cuda-memcheck ./t1218 
========= CUDA-MEMCHECK 
Reduced sum of Array elements = 523776 
Value should be: 523776 
========= ERROR SUMMARY: 0 errors 
$ 

Примечания:

1. Ядра и ваш код, как написано зависят от N, являющегося точным кратным размеру блока threadblock (256). Это удовлетворено для этого случая, но все будет нарушено, если это не так.

2. Я не вижу никаких доказательств proper cuda error checking. Здесь ничего бы не получилось, но его хорошая практика. В качестве быстрого теста запустите свой код с cuda-memcheck, как я это сделал.