[CUDA编程]基础入门例程4

例程地址：https://github.com/TycoonL/awesome-cuda.git/

7.basic_reduce

本例程实现用全局内存进行数组规约（N<128）。

注意：对于多线程程序，程序的执行进程是不同步的，会造成结果错误，所以在每次规约后需要用‘__syncthreads()’同步线程。

__global__ void reduce(float* in, float* out, int N) {
	int tid = threadIdx.x;
	float *x = in + blockIdx.x * blockDim.x;
	__syncthreads();
	for (int i = blockDim.x / 2; i > 0; i /= 2) {
		if (tid < i) {
			x[tid] += x[tid + i];
		}
		__syncthreads();
	}
	if (tid == 0) {
		out[blockIdx.x] = x[0];
	}
}
//N=100,result:5050.000000,Execution time: 0.018432 ms

8.basic_reduce_share

本例程是对例程7的改进，用共享内存进行数组规约。

共享内存是一种能被程序员直接操控的缓存，作用：

1. 减少核函数对全局内存的访问次数，实现高效的线程块内部通信
2. 提高全局内存的合并度

核函数：

__global__ void reduce(float* in, float* out, int N) {
	int tid = threadIdx.x;
	int n = tid + blockIdx.x * blockDim.x;
	__shared__ float s_y[128];//声明静态共享内存
	s_y[tid] = (n < N) ? in[n] : 0.0;
	for (int i = blockDim.x / 2; i > 0; i /= 2) {
		if (tid < i) {
			s_y[tid] += s_y[tid + i];
		}
		__syncthreads();
	}
	if (tid == 0) {
		out[blockIdx.x] = s_y[0];
	}
}
//N=100,result:
//加上__syncthreads()：5050.000000,Execution time: 0.022528 ms
//去掉__syncthreads()：5050.000000,Execution time: 0.015360 ms

数组规约使用共享内存还有两个好处：

1. 不再要求全局内存数组的长度N为线程块大小的整数倍
2. 在规约过程中不会改变全局内存数组中的数据

9.basic_reduce_share_dynamic

本例程将例程8的静态内存修改成动态共享内存

动态的共享内存能防止长度定义错误

做以下两处修改：

reduce<<<grid_size, block_size, N*sizeof(float)>>>(d_in, d_out, N);

extern __shared__ float s_y[];