CUDA精讲（2）-- GPU体系结构之线程执行

GPU体系结构

本讲是CUDA精讲的第二部分，在CUDA精讲（1）中主要列出了CUDA编程的一些基本概念。为了进一步地深入CUDA的系统优化我们需要了解GPU的硬件体体系结构。大部分的处理器体系结构都可以分成计算、存储、控制三部分。GPU中主要强调计算（Thread exection）与存储（Memory hierarchy）两部分，下面就以下两部分进行展开。

• 线程执行（Thread execution）
• 内存层次结构（Memory hierarchy）

Thread execution

1. CUDA线程执行特点：

• thread组以blocks的形式体现，block对应硬件中的SM。
• SM中有shared memory，同一个block中的线程可以借助block进行通信
• block中的所有线程全部执行完毕之后，scheduler才能调度下一个线程。因此有可能出现一个线程执行慢导致整个block执行慢。

2. 线程调度的基本单位Warp

• 一个warp包含32个线程
• 在任意时刻，在SM中只有一个warp scheduler在执行。
• 一个warp中的线程组必须执行相同的指令

3. Warp Divergence
   当执行的kernel函数中存在if分支的时候比如：

if(foo(threadIdx.x)){
	do_A();
} else {
	do_B();
}

由于warp中的线程必须执行相同的指令，因此这些线程会限制性函数A，然后再执行函数B，最后再根据条件将函数A和函数B的结果进行合并。执行过程如下图所示：

因此if中的分支会导致warp中的thread串行执行分支中的代码。

• 避免warp diverging的方法：
  分支的判断粒度是warp的整数倍，例如：

if(threadIdx.x / WARP_SIZE > 2) {...}

4. iteration Divergence, 迭代发散，当一个kernel中存在一个迭代循环的时候，有可能因为一个thread的循环次数过多而导致一个warp中的所有线程都要等待。

5. 通过unroll减少for循环中条件判断的次数：
   例如：
   
   还可以通过一个简单的方法来表示：
   
   编译器会自动的将这个语法展开。

6. 原子操作，当多个线程同时去更改一个数据的时候，可以通过原子操作进行同步，常见的原子操作指令有：

   • 加减， atomicAdd
   • max，min atomicMax
   • 自增，自减