reduccionParalela.cu

#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#include <cuda_runtime.h>
#include <math.h>

#define HILOS 45

/**
 * Sumatorio de los valores de un array mediante paralelismo en NVIDIA 
 */
__global__ void reduccionParalela(int array[]){
	int myID = threadIdx.x;

	if(HILOS%2 != 0){
		if(myID == (HILOS-1)){
			array[0] += array[myID];
		}
	}
	int salto = HILOS/2;
	
	__syncthreads();

	while(salto)
	{
		// Solo trabajan la mitad de los hilos
		if(myID < salto){

			array[myID] = array[myID] + array[myID+salto];
		}
		__syncthreads();
		//Comprobamos que la mitad no sale impar y si es así se suma el valor perdido
		if(salto % 2 != 0 && salto!=1){
			if(myID == salto-1){
				array[0] = array[0] + array[myID];
			}
		}
		__syncthreads();
		salto = salto/2;
	}
	//El valor ahora en array[0] es el doble exacto al que queremos, lo partimos y ya. //Solo funciona con enteros?
}

int main(int argc, char** argv){

	srand(time(NULL));

	int *hst_array = (int*) malloc(sizeof(int)*HILOS);
	int *dev_array;
	int total=0;
	cudaMalloc((void**)&dev_array,sizeof(int)*HILOS);
	for(int i = 0; i<HILOS;++i){
		hst_array[i]=i;
		printf("%d ",hst_array[i]);
		total += hst_array[i];
	}
	cudaMemcpy(dev_array,hst_array,sizeof(int)*HILOS,cudaMemcpyHostToDevice);
	reduccionParalela<<<1,HILOS>>> (dev_array);
	cudaMemcpy(hst_array,dev_array,sizeof(int)*HILOS,cudaMemcpyDeviceToHost);
	printf("\nReducción: %d - Bucle: %d\n",hst_array[0],total);


	return 0;
}