最近在研究并行运算的规约算法,在看《GPU高性能编程CUDA实战》这本书中点积运算时,有些问题想了很久,记录下来;
注点积公式:(dot(A,B)=a1*b1+a2*b2+...+an*bn)
书上例子算点积运算时分为了以下几步:
1.申请共享内存cache;
__shared__ float cache[threadsPerBlock]
这里要注意一个概念:一旦这样声明共享内存,就会创建与线程块的数量相同的数组cache,即每个线程块都会对应一个这样的数组cache且不同线程块中的共享内存是无法交流的;
在这个例子中,共享内存的大小与每个线程块中的线程个数相同;
2.将每个线程块中每个线程计算的元素(a1*b1形式)加起来并放入共享内存;
int tid = threadIdx.x + blockIdx.x * blockDim.x;
int cacheIndex = threadIdx.x;
float temp = 0;
while (tid < N)//N为数组元素个数
{
temp += a[tid] * b[tid];
tid += blockDim.x * gridDim.x;
}
cache[cacheIndex]=temp;
__syncthreads();
该例线程块和线程都申请的是一维的,所以tid就代表所有线程块中的所有线程数量;
在while循环这困扰了我很久,不明白为什么需要求和?直