// --Kernel-Funktion fuer die Parallele Reduktion

__global__ void parallel_reduce(int* arr, int radix2_size,int nsteps) {
// --Wer bin ich?
int index=(blockDim.x*blockIdx.x+threadIdx.x)*2,s=0;

for(int i=0;i<nsteps;i++) {
    s=1<<i;		// --s: Schrittweite
    // --Ueberpruefe, ob Thread aktiv rechnen muss.
    if (threadIdx.x % s == 0)
        arr[index]+=arr[index+s];

    // --Synchronisiere mit allen anderen Threads im Block
    __syncthreads();
    }

}



// --Hier werden die Gitter- und Blockdimensionen definiert und der Kernel gestartet.

const int radix2Threads = 9;		// --2^.. Threads

extern "C" void reduce(int* arr, int radix2_size) {

// --Dimension eines eindimensionalen Blocks
dim3 dimBlock(1<<radix2Threads);

// --Berechnung der Anzahl der Bloecke im eindimensionalen Gitter
dim3 dimGrid(1<<(radix2_size-radix2Threads-1));

// --Berechnung der Anzahl der Reduktionsschritte pro Block
int nreduction=radix2Threads+1;

// --Starte die Kernelfunktion auf dem Gitter
parallel_reduce<<<dimGrid,dimBlock>>>(arr,radix2_size,nreduction);

}