CUDA __shared__ thread、block、grid之间的一维关系 （例子chapter5 dot点积（GPU高性能编程））

#include <iostream>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "book.h"
#include "cpu_bitmap.h"

//1.minFun  对于每个Grid而言，选取block数比较少的，这样可以减少每个grid中最后一个block内thread的冗余。

int imin(int a,int b)
{
	if(a>b)
	{
		return b;
	}else{
		return a;
	}
}

const int N = 33*1;//整体线程的数量，即向量（a）*(b)中每个向量数组的大小
const int ThreadNumsPerBlock =256;//每个block中thread的数量
const int BlockNumsPerGrid = imin(32,(N+ThreadNumsPerBlock-1)/ThreadNumsPerBlock);//每个grid中block的数量，

__global__ void dot(float* a,float* b,float* c)//输入两个向量：a，b ；结果c大小为一个grid内block的大小，grid内每个block的位置对应数组c每个下标对应的数据，每个数据为多个grid相同的block下标下的thread的和，如图-1。
{
   int	everyThreadIndex  =  threadIdx.x+blockIdx.x*blockDim.x;//获取一个grid中thread的index
   int  cacheIndex = threadIdx.x;//获取一个block中thread的index
   __shared__ float cache[ThreadNumsPerBlock];

   float temp = 0;
   while(everyThreadIndex<N)//多个grid的跳跃index，
   {
	   temp += a[everyThreadIndex]*b[everyThreadIndex];
	   everyThreadIndex += blockDim.x*gridDim.x;//
   }
   cache[cacheIndex] = temp;//得到多个grid中 一个block里相应的一个thread的乘积的和
   __syncthreads();

   int i= blockDim.x/2;
   while(i!=0)   //一直纳闷这个怎么在多个线程中执行啊？应该是：一个thread执行到这里后，__syncthreads()之前一个语句等啊等，等到所有一个block里的threads全都执行一遍后，才往下进行
   {
	   if(cacheIndex<i)
	   {
		   cache[cacheIndex]= cache[cacheIndex]+cache[cacheIndex+i];//while里的这句是在block里不同的thread里执行的
	   }
	   __syncthreads();
	   i=i/2;//当每个thread都执行后，才能进行长度折半
   }
   if(0==cacheIndex)
   {
     c[blockIdx.x] = cache[0];
   }
}

int main(void)//
{
	float  *host_a,*host_b,*host_c;
	float  *device_a,*device_b,*device_c;

	host_a = (float*)malloc(N*sizeof(float));
	host_b = (float*)malloc(N*sizeof(float));
	host_c = (float*)malloc(BlockNumsPerGrid*sizeof(float));

	  for (int i=0; i<N; i++) {
        host_a[i] = i;
        host_b[i] = i*2;
    }

	//gpu上分配内存
	cudaMalloc((void**)&device_a,N*sizeof(float));
	cudaMalloc((void**)&device_b,N*sizeof(float));
	cudaMalloc((void**)&device_c,BlockNumsPerGrid*sizeof(float));

	cudaMemcpy(device_a,host_a,N*sizeof(float),cudaMemcpyHostToDevice);
	cudaMemcpy(device_b,host_b,N*sizeof(float),cudaMemcpyHostToDevice);

	dot<<<BlockNumsPerGrid,ThreadNumsPerBlock>>>(device_a,device_b,device_c);

	cudaMemcpy(host_c,device_c,BlockNumsPerGrid*sizeof(float),cudaMemcpyDeviceToHost);
    
	float c=0;
	for(int i=0;i<BlockNumsPerGrid;i++)
	{
		c = c + host_c[i];
	}
	  #define sum_squares(x)  (x*(x+1)*(2*x+1)/6)
	std::cout<<"结果1："<<c<<"结果2："<<2 * sum_squares( (float)(N - 1) )<<std::endl;

	cudaFree(device_a);
	cudaFree(device_b);
	cudaFree(device_c);

	free(host_a);
	free(host_b);
	free(host_c);

	return 0;
}

while(i!=0)
   {
	   if(cacheIndex<i)
	   {
		   cache[cacheIndex]= cache[cacheIndex]+cache[cacheIndex+i];//while里的这句是在block里不同的thread里执行的

                   __syncthreads();
            }

         i=i/2;//当每个thread都执行后，才能进行长度折半 

   }