阅读 《大规模并行处理器程序设计》影印版心得 第四章 CUDA Threads

 

4.1 CUDA  Thread Organization

 

具体例子：一个grid中有N个block，但是以一维的形式组织起来。每一个block中有M个线程，也以一维的形式组织起来。则任何一个block中的线程可以号可以用公式 threadID = blockIdx.x *blockDim.x +threadIdx.x来计算。

 

两个变量：gridDim和blockDim, gridDim.x, gridDim.y

 

注意：grid中的block用二维方式来组织，block中的thread用三维形式来组织。每一个block中最多有512个线程thread

 

 

4.2 Using blockIdx and threadIdx

 

通过tile技术（对数据的细分）增加线程数量，使用多block创建更多线程thread

 

这里的一个难点是如何通过blockIdx和threadIdx来计算应该取两个相乘矩阵的哪一行和哪一列。图4-3应该有很好的解释作用。这里的关键是grid中各个block的排列关系和每个block中各个thread的排列关系，图4.2中是标准的排列方式。 -- 最好测试一下，换一种排列关系的效果？ 正确与否？ 如果正确，效率是否有区别？

 

 

4.3 Synchronization and Transparent Scalability

 

这一块比较简单。主要是利用函数_syncthreads()来实现同一个block的不同线程之间的同步。需要注意的是，if-then-else中的_syncthread()函数如果两个分支都有，则两个分支中的算两个同步点。

 

还有一个，就是因为有同步的可能，则同一个block中的thread将被分配同样的资源量。因此，系统同时能够运行的block的个数就跟资源挂钩了。当资源很多时，同时运行的block个数就可以很多；反之，就可以很少。 貌似这个多少不由程序员控制，显卡自己能控制。

 

4.4 Thread Assignment

 

SM streaming multiprocessor  在GPU中，执行资源被组织为SMs

 

在GT200芯片显卡中，最多可以有30个SM，每个SM最多可以容纳8个block，每一个SM最多可以容纳1024个threads。 哪个先到算哪个。

 

SP和SM的关系？

 

4.5 Thread Scheduling and Latency Tolerance

 

在GT 200显卡中，当一个block被放到一个SM上运行时，该block又被分为32个线程一组的warps. warp是SM中线程调度的基本单位。

 

latency hiding  调度非等待的warp进入执行态，使SM始终处于忙状态，提高了效率，降低了warp读取global memory时的延迟开销。主要是warp的调度是零费时的，不会引入任何额外开销

 

因此，为了性能优化，我们可以考虑增加一个SM中的线程数，因此增加其warp数量。 这样，可供调度的warp增加了，SM忙的概率就将增加。