CUDA-CODE6-共享和同步

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <iostream>
#include <cuda.h>
#include"device_functions.h"
/*
线程块Block的数量限制是：65535
我的机器上maxThreadPerBlock = 1024.
也就是一个Block中最多1024个线程
---
本篇主要是共享变量的使用
*/

const int  N  = (33 * 1024);

#define imin(a,b) ( a < b ? a:b )

const int threadPerBlock = 512;
const int blockPerGrid = imin(32 , (N + threadPerBlock -1 ) / threadPerBlock );

__global__ void  dot(float *a , float *b , float* c)
{
	__shared__ float cache[threadPerBlock];
	int tid = threadIdx.x + blockIdx.x * blockDim.x;//获得线程索引
	//			printf("tid = %d\n",tid);
	int cacheIndex = threadIdx.x;
	float temp = 0;
	while( tid < N)	//必须检查是否在正确的偏移范围类内
	{
		temp += a[ tid ] * b[ tid];
		tid += blockDim.x * gridDim.x;  
	}
	//设置cache中相应位置上的值
	cache[ cacheIndex ] = temp;
	//对线程块中的线程进行同步
	__syncthreads();//每个线程都必须保证前面的操作执行完了之后才能执行下面的语句，
	//也就是说，如果有的线程提前完成了前面的操作，而其他线程还没完成，那么该线程将进入挂起状态
	//这其实是一种约定，大家都必须先把前面的X干完，才能一起进入下个阶段
	//必须要这样保证，因为如果M提前进入下一个阶段的话，他将去取其他cache位置上的数，但是他会发现，有些数字根本还没装进来，因为其他线程还没计算完，如果他执意要进行工作，那么必定会导致最终的计算错误。所以，这里必须要同步，要大伙儿一起进入下一个阶段才能保证结果的正确。
	//归约运算
	int i = blockDim.x / 2;
	while( i != 0)
	{
		if( cacheIndex < i )
			cache[ cacheIndex ] += cache[ cacheIndex + i ];
		__syncthreads();//和上面的解释相同道理，因为都是读cache数据，必须要保证所有线程的写操作都完成，才能继续下一轮的读取，再写，再读...
		i /= 2;
	}
	if( cacheIndex == 0)
		c[ blockIdx.x] = cache[0]; 
}

int main()
{
	float *a , *b , c , *partial_c;
	float *dev_a, *dev_b , *dev_partial_c;
	//在CPU上分配内存
	a = (float*)malloc( N * sizeof(float));
	b =  (float*)malloc( N * sizeof(float));
	partial_c =  (float*)malloc( blockPerGrid * sizeof(float));
	//在GPU上分配内存
	cudaMalloc( (void**)&dev_a , N * sizeof(float ) );
	cudaMalloc( (void**)&dev_b , N * sizeof(float ) );
	cudaMalloc( (void**)&dev_partial_c , blockPerGrid * sizeof(float ) );

	//在CPU上为数组a,b赋初值
	int i = 0;
	for( i = 0; i < N; i++)
	{
		
		*(a+i) = i ;	 
		*(b+i) = i ;
		//a[i] = i;
		//b[i] = i;
		//printf("%f\n",*(a+i));
		 
	}
	//将数组a, b 复制到GPU上去计算
	cudaMemcpy( dev_a , a , N * sizeof( float ) , cudaMemcpyHostToDevice);
	cudaMemcpy( dev_b , b , N * sizeof( float ) , cudaMemcpyHostToDevice);

	dot<<<blockPerGrid,threadPerBlock>>>( dev_a , dev_b , dev_partial_c );//那么将会有128 * 128 个核函数运行
	 
	cudaMemcpy( partial_c , dev_partial_c , blockPerGrid * sizeof( float ) , cudaMemcpyDeviceToHost);
	 //在CPU上完成最终的求和运算
	c = 0;
	for(int i = 0 ; i < blockPerGrid ; i++)
	{
		c += *(partial_c + i );
		
	}
	
	cudaFree( dev_a );
	cudaFree( dev_b );
	cudaFree( dev_partial_c );
	//释放CPU上的内存
	free( a );
	free( b );
	free( partial_c );
	printf("c = %f" , c);
	getchar();
	return 0;
}