文章目录

线程协作实例: <数组和>

__global__ void sum(float *a)            
	__shared__ float cache[threadsPerBlock];
	int tid = threadIdx.x + blockIdx.x * blockDim.x;
	int cacheIndex = threadIdx.x;
	/*
		generate the value of cache[];
	*/
	int i = blockDim.x/2;
	while( i != 0 ){
		if(cacheIndex < i)
			cache[cacheIndex] += cache [cacheIndex + i ];
		__syncthreads();
		i /= 2;
	}
	
	if (cacheIndex == 0)
		c[blockIdx.x] = cache[0];
}

A. 共享内存方式：

• 关键字share float cache[threadsPerBlock]; //相当于在标准C中声明一个static类型的变量
• 对于GPU上启动的每个线程块，CUDA C编译器都将创建该变量的一个副本。线程块中的每个线程都共享这块线程。但线程却无法看到也不能修改其他线程块的变量副本。
• 共享内存缓冲区驻留在物理GPU上，而不是系统内存；大大降低了访问延迟。
• syncthreads(); 用以保证线程块中的每个线程：都执行完__syncthreads()前面的语句, 才进行后面的语句。

B. 代码基本思想：
每个线程将cache[]中的两个值相加起来，然后将结果保存回cache[]。由于每个线程都将两个值合并为一个值，那么在完成这个步骤后，得到的结果数量就是计算开始数值数量的一半，如此重复。

C. 为什么只有cacheIndex == 0的线程执行保存操作？

• 因为只有一个值需要写入内存，因此只需要一个线程来执行这个操作。

D. 思路总结：

• 为输入数组和输出数组分配主机内存和设备内存
• 填充输入数组a[],b[],然后通过cudaMemcpy()将它们复制到设备上。
• 在调用计算数组和的核函数时指定线程格中线程块的数量以及每个线程块中的线程数量。

E. 线程块数量计算分析：
有N个数据元素，那么通常只需要N个线程来计算和。但在这里的情况中，线程数量应该为threadsPerBlock的最小整数倍，并且要>=N .
因此建议启动线程块数量 == (N + (threadsPerBlock-1))/threadsPerBlock
即：
dim3 dimBlock(8,8,1);
dim3 dimGrid((N1+dimBlock.x-1)/dimBlock.x, (N2+dimBlock.y-1)/dimBlock.y,1);