CUDA global_reduce实现

CUDA global_reduce实现

#include <stdio.h>
#include <time.h>
#include <stdlib.h>
#include <cuda_runtime.h>
__global__ void global_reduce(float* d_out, float* d_in)
{
	int myID = blockIdx.x * blockDim.x + threadIdx.x;//得到该线程在全局的线程号
	int idx = threadIdx.x;//得到该线程在所在块的线程号
	for (unsigned int s = blockDim.x / 2; s > 0; s=s/2)
	//算法核心：将一个块中的线程分成2部分，线程块左边的线程操作自己线程所控制的内存再加上与自己操作的内存相差s个相对单位的内存
	//即d_in[myID] += d_in[myID + s]，而右边的线程不做任何操作，之后又将左边的线程分为2部分重复上述操作直到s=0。
	{
		if (idx < s)
		{
			d_in[myID] += d_in[myID + s];
		}
		__syncthreads();//让线程等待，避免下次循环出错
	}
	if (idx == 0)//按照上面的模式，每一个块的最终结果都会记录在该块中第一个线程所操作的值里。
	{
		d_out[blockIdx.x] = d_in[myID];//将结果存入输出数组里
	}
}


int main()
{
	float h_in[1024];
	int i;
	for (i = 0; i < 1024;i++)
	{
		h_in[i] = i;//初始化
	}
	float h_out[1];
	
	float* d_in;
	float* d_out;
	//分配GPU内存
	cudaMalloc(&d_in, sizeof(float) * 1024);
	cudaMalloc(&d_out, sizeof(float));
	//将cpu上的值拷贝到GPU上
	cudaMemcpy(d_in, h_in, sizeof(float) * 1024, cudaMemcpyHostToDevice);
	//启动kernal，此处设置了1个block和1024条thread
	global_reduce << <1, 1024 >> > (d_out, d_in);
	//GPU上的值拷贝回CPU
	cudaMemcpy(h_out, d_out, sizeof(float), cudaMemcpyDeviceToHost);
	//检查并行结果与串行结果是否一致
	float sum = 0;
	for (int i = 0; i < 1024; i++)
	{
		sum += h_in[i];
	}
	printf("%f\n", h_out[0]);
	printf("%f", sum);
	return 0;
}