（1）线程束究竟是什么？

概念上讲，线程束是SM用SIMD方式所同时处理的工作粒度。当一个线程块被指定到一个SM上，线程块中的所有线程被分成了线程束，所以优化工作负载要考虑到线程束。或者这样说，我们从软体角度看线程，线程是独立的且并行执行，对硬体来说，却是被组织成32个线程为一堆的线程束来执行指令。因此，逻辑上我们每个线程块可以是一维，二维，三维线程，但是物理上都是以32个线程为单位的一维物理布局。

杰伦：线程块中线程数应该是32的倍数，不然最后那个没有被填满的线程束，会和其他填满的线程束一样消耗SM的资源，比如宝贵的寄存器。

（2）条件分支与线程束的关系

有一点非常重要，在一个线程束中所有线程在每个周期必须执行相同的指令，也就是说一个线程束中的所有线程，在面临条件分支的时候，集体选择同一条件执行，如果该线程束中有16个线程是选择a条件执行，另外16个是选择b条件执行，则实际的工作情况是：16个选择a在执行的时候将会禁用另外16个线程执行b，也就是串列执行不同的条件分支，这个对性能的影响是很恶劣的。举个栗子：

__global__ void kernel(float *c)
{
int index = threadIdx.x + blockIdx.x * blockDim.x;
float a =0,b=0;
if(index%2==0)
a = 100.0f;
else
b = 200.0f;
c[index] = a+b;
}

改进版（以线程束为单位，不再是以线程IDX奇偶，而是线程束的奇偶来分）

__global__ void kernel(float *c)
{
int index = threadIdx.x + blockIdx.x * blockDim.x;
float a =0,b=0;
if(（index/wrapSize)%2==0)//线程束为偶数执行，这样可行的先决条件是，有一半线程执行a，有一半线程执行b
a = 100.0f;
else
b = 200.0f;
c[index] = a+b;
}

杰伦：在内核函数中，尽量避免数据分化，解决方法就是在条件分支处以线程束为单位。

（3）增大并行性来提高性能

实验表明，当我们的blockDim.x是32的倍数，能获增加载入吞吐量，而块数较多时（当然不能超过硬体限制），可以有更多活跃的线程束，这直接影响著可实现占用率，从而影响性能。但是，并不是说可实现占用率越高，性能就越好，在实验的时候要注意调整块大小。

杰伦：当我们想通过设置块的大小，网格的大小来获得最佳性能时，不能单单考虑到可实现占用率，载入吞吐量，要在指标之间寻求一个好的平衡，原则就是，块数偏多时会增加并行性，块第一维（第一维调整对载入吞吐量有绝对的影响）一定要是32的倍数。

（4）CUDA内存模型

无论是CPU还是GPU都遵循时间局部性和空间局部性，这就使得CPU和GPU内存结构设计都是低延迟低容量的层次结构，如cpu中（寄存器、缓存、主存、磁碟存储器，容量越大，延迟越高）。在CUDA的内存层次中，排序为（寄存器，共享内存，本地内存，常量内存，纹理内存，全局内存），如图所示：

GPU内存结构
CPU内存结构
  1. 寄存器：低容量低延迟，线程级的内存，为一个线程私有，生命周期与核函数一样，核函数中声明的没有其他修饰符的自变数，通常是存储在寄存器中，值得注意的是，在核函数中使用较少的寄存器将使在SM上有更多的常驻线程块，每个SM上并发线程块越多，使用率和性能越高，所以这也是为什么要避免半个线程束这种不活跃的线程束的存在，原因就在于浪费了SM的寄存器。
  2. 本地内存：寄存器的备胎，寄存器中变数的溢出地，有限定条件的自变数（没理解到这是什么变数）但是当作为寄存器数据的溢出内存用的时候，却与全局内存在一个地方，于是高延迟低带宽。一般来说，对于计算能力在2.0及以上的GPU来说，本地内存是SM的一级缓存，设备的二级缓存。
  3. 共享内存：标记：__share__ 高带宽低延迟，线程块级，类似于CPU中的一级缓存，但他可编程。生命周期与线程块一致，但需注意，片上内存64KB中一部分分配给共享内存，一部分分配给GPU一级缓存（可通过cudaFuncSetCacheConfig()手动配置），我的共享内存是48KB。共享内存是线程间通讯的必要内存，但是要注意搭配线程同步。在《CUDA by Example》--chapter05 code中有讲过共享内存的具体用法。
  4. 常量内存：标记：__constant__ ，在全局空间和核函数之外进行声明，对所有线程可见，初始化：cudaMemcpyToSymbol(); 适用范围：当所有线程从相同的内存地址中读取数据时，常量内存表现良好。具体用法，见书《CUDA by Example》--chapter06 。
  5. 纹理内存：。。。具体使用还是看《CUDA by Example》--chapter07吧！
  6. 全局内存：用得最多，实际意义上的显存，也就预示著它是高容量高延迟。我的显卡是N卡1060，显示全局内存为2G，我们使用cudaMalloc()与cudaFree()就是在动态得分配全局内存。由于不存在块间同步，因此不同线程块并发得修改全局内存中的同一位置可以会出现未定义行为。
  7. 注意，纹理，常量，全局内存都是可以被所有线程访问，且都是在主机上配置，但是纹理和常量是只读访问。

（5）

推荐阅读：

相关文章