So finden Sie die Summe der Array in CUDA durch Reduktion

Question

Ich bin eine Funktion implementieren, um die Summe eines Arrays mit Reduktion zu finden, mein Array haben 32 * 32 Elemente und seine Werte ist 0 ... 1023. Die mein erwarteter Summenwert ist 523776, aber mein Ergebnis ist 15872, es ist falsch. Hier ist mein Code:

#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

Es gibt mindestens zwei Probleme im Code

1. Sie sind in Ihrem dev_b Array atomicAdd auf das erste Element zu tun, aber Sie initialisieren nicht, dieses Element zu ein bekannter Wert (dh 0). Sicher, bevor Sie den Kernel ausführen, kopieren Sie b zu dev_b, aber da Sie b nicht zu irgendwelchen bekannten Werten initialisiert haben, wird das nicht helfen. Das Array b wird in C oder C++ nicht automatisch auf Null initialisiert, wenn Sie daran gedacht haben. Wir können dies beheben, indem wir b[0] auf Null setzen, bevor wir b zu dev_b kopieren.

Ihr Reduzierungskern ist so geschrieben, dass er einen 1D-Fall behandelt (d. H. Der einzige verwendete Threadindex ist ein 1D-Threadindex basierend auf den Werten .x), Sie starten jedoch einen Kernel mit 2D-Threadblocks und -Grids. Dieser Konflikt wird nicht richtig funktionieren und wir müssen entweder einen 1D-Threadblock und ein Gitter starten oder den Kernel neu schreiben, um mit 2D-Indizes zu arbeiten (d. H. .x und .y). Ich habe die ehemalige (1D) gewählt.

hier ein ausgearbeitetes Beispiel mit den Änderungen an Ihrem Code ist, so scheint es, das richtige Ergebnis zu produzieren:

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

Hinweise:

1. Der Kernel und Ihren Code geschrieben abhängig von N ist ein genaues Vielfaches der Threadblockgröße (256). Das ist für diesen Fall erfüllt, aber die Dinge werden brechen, wenn es nicht so ist.

2. Ich sehe keinen Beweis von proper cuda error checking. Es hätte hier nichts gefunden, aber es ist eine gute Übung. Als Schnelltest starte deinen Code mit cuda-memcheck wie ich es hier gemacht habe.