目录

一、初步认识

二、Warps and Thread Blocks

三、Warp Divergence 

四、资源分区

五、延迟隐藏

六、占有率

七、Synchronize

---

一、初步认识

warp 是 SM 的基本执行单元。一个 warp 包含 32 个并行 thread，这32个 thread 执行于 SMIT模式。即所有 thread 执行同一条指令，并且每个 thread 会使用各自的 data 执行该指令

1. 线程调度：GPU 的调度单元以 warp 为单位进行调度，而不是单个线程。即整个 warp 被分配到一个流多处理器（SM）上并一起执行
2. 分支处理：若 warp 中的所有线程都采取相同的分支路径（如：都满足某个条件语句），则其会继续同步执行。但是，若线程在分支上有不同的路径（即分歧），则 warp 会执行每个路径，但不是所有线程都会在每个路径上活跃。这可能导致效率下降，因为即使某些线程在特定路径上没有工作，整个 warp 也必须等待该路径完成
3. 性能优化：了解 warp 的行为对于优化 CUDA 程序的性能至关重要。如：为了确保高效执行，开发人员可能需要确保他们的代码减少 warp 分歧。
4. 占用率：在 CUDA 中，占用率是一个重要的性能指标，表示每个 SM 上激活的 warps 与 SM 可以支持的最大 warp 数量的比例。更高的占用率通常意味着更好的硬件利用率

即 warp 是 NVIDIA GPU 中执行并行操作的基本单位

二、Warps and Thread Blocks

block可以是一维二维或者三维的，但是，从硬件角度看，所有的thread都被组织成一维，每个thread都有个唯一的index

每个block的warp数量可以由下面的公式计算获得：

一个warp中的线程必然在同一个block中，若block所含线程数目不是warp大小的整数倍，即使凑不够warp整数倍的thread，硬件也会为warp凑足，只不过那些thread是闲置状态。注意：即使这部分thread是闲置的，也会消耗SM资源 

三、Warp Divergence 

同一个 warp 中的 thread 必须执行相同的指令，若这些线程在遇到控制流语句时，进入了不同的分支，那么同一时刻除了正在执行的分支外，其余分支都会被阻塞，十分影响性能。如下图： 

当代码较为简单时，GPU会自动进行代码优化，即进行分支预测

在GPU编程中，预测变量是一个用于指示分支执行路径的变量，其值为1或0。当条件状态少于某个阈值时，编译器会将一个分支指令替换为预测指令。若预测变量为1，则执行预测为真的分支；若预测变量为0，则执行预测为假的分支

当编译器检测到某个条件判断语句的执行频率较低时，其会假设这个条件判断的结果在大多数情况下是已知的，从而将其替换为预测指令。这样一来，程序在运行时就不需要每次都进行实际的条件判断，而是直接根据预测结果执行相应的代码路径

虽然分支预测可以在一定程度上提高程序的执行效率，但并不总是准确的。当预测错误时，程序需要回滚到正确的分支路径上继续执行，这会带来一定的性能开销

四、资源分区

一个 warp 的 context 包括以下三部分：

1. Program counter（程序计数器）
2. Register（寄存器）
3. Shared memory（共享内存）

在同一个 执行context 中切换是没有消耗的，因为在整个 warp 的生命期内，SM 处理的每个 warp 的 执行context 都是 on-chip 的

每个 SM 有一个 32位 register 集合放在 register file 中，还有固定数量的 shared memory，这些资源都被 thread 瓜分了。由于资源是有限的，若 thread 较多，那么每个 thread 占用资源就较少，thread 较少，占用资源就较多，这需要根据需求作出平衡

资源限制了驻留在 SM 中 blcok 的数量，不同的device，register 和 shared memory 的数量也不同，就像上面介绍的 Fermi 和 Kepler 的差别。若没有足够的资源，kernel 的启动就会失败

当一个block或得到足够的资源时，就成为 active block。block 中的 warp 就称为 active warp。active warp 又可以被分为下面三类：

1. Selected warp（被选中的）
2. Stalled warp（停滞的）
3. Eligible warp（符合条件的）

SM 中 warp调度器 每个 cycle 会挑选 active warp 送去执行，一个被选中的 warp 被称为 selected warp，没被选中，但是已经做好准备被执行的称为 Eligible warp，没准备好要执行的称为 Stalled warp。warp 适合执行需要满足下面两个条件: 

1. 32个CUDA core有空
2. 所有当前指令的参数都准备就绪

Kepler 任何时刻的 active warp 数目必须小于等于64个。selected warp 数目必须小于等于4个。若一个 warp 阻塞了，调度器会挑选一个 Eligible warp 准备去执行

CUDA编程中应重视对计算资源的分配：这些资源限制了 active warp 的数量。因此，必须掌握硬件的一些限制，为了最大化GPU利用率，须最大化 active warp 的数目

五、延迟隐藏

指令从开始到结束消耗的 clock cycle 称为指令的 latency。当每个 cycle 都有 eligible warp 被调度时，计算资源就会得到充分利用，基于此就可以将每个指令的 latency 隐藏于 issue 其它 warp 的指令的过程中（issue 指令是指将指令发送到执行单元进行执行的操作）

和CPU编程相比，latency hiding 对 GPU 非常重要。CPU cores 被设计成可以最小化一到两个 thread 的 latency，但是GPU的 thread 数目可不是一个两个那么简单

当涉及到指令latency时，指令可以被区分为下面两种：

1. Arithmetic instruction（算术指令）
2. Memory instruction（内存指令）

Arithmetic instruction latency是一个算数操作的始末间隔。另一个则是指 load 或 store 的始末间隔。二者的 latency 大约为：

1. 10 - 20 cycle for arithmetic operations（10 - 20 个算术运算周期）
2. 400 - 800 cycles for global memory accesses（400 - 800 次全局内存访问）

下图是一个简单执行流程，当 warp0 阻塞时，执行其他的 warp，变为 eligible 时重新执行

如何评估 active warps 的数量来hide latency呢？Little’s Law提供一个合理的估计：

所需warps的数量 = 延迟 * 吞吐量

 Arithmetic operations

对于 Arithmetic operations 来说，并行性可以表达为用来 hide  Arithmetic latency 的操作的数目

Throughput 为 32，即每个 SM 在每个时钟周期能够执行 32 个操作

latency 为 20，这即一个操作需要 20 个时钟周期才能完成

由于 latency 为 20，在这 20 个时钟周期内，SM 要保持吞吐量为 32 的操作执行效率。那么在这 20 个时钟周期内总共需要的操作数是  个（因为每个周期要执行 32 个操作，一共 20 个周期）。又因为每个 warp 能提供 32 个操作（每个 warp 执行同一种指令，对应 32 个操作），所以计算得出每个 SM 需要  个 warp，这样才能在 latency 期间，保证每个 SM 都有足够的操作来维持吞吐量，从而充分利用计算资源

Memory operations

对于Memory operations，并行性可以表达为每个cycle的byte数目

因为memory throughput（带宽）总是以GB/Sec为单位，需要先作相应的转化。可以通过下面的指令来查看device的memory frequency（内存频率）：

nvidia-smi -a -q -d CLOCK | fgrep -A 3 "Max Clocks" | fgrep "Memory"

以Fermi为例，其memory frequency可能是1.566GHz，Kepler的是1.6GHz。那么转化过程为：

800乘上92约等于上图中的74，这个数字是针对整个device的，而不是SM

以Fermi为例，假设每个thread的任务是将一个float（4 bytes）类型的数据从 global memory 移至 SM 用来计算，应需要大约18500个thread，即579个warp来隐藏所有的memory latency

Fermi 有16个SM，所以每个SM需要 579/16=36 个warp来隐藏 memory latency

六、占有率

当一个 warp 阻塞了，SM 会执行另一个 eligible warp。 理想情况是，每时每刻到保证 cores 被占用。 Occupancy 就是每个 SM 的 active warp 占最大 warp 数目的比例：

grid 和 block 的配置准则：

• 保证 block 中 thrad 数目是32的倍数
• 避免block太小：每个blcok最少 128 或 256 个thread
• 根据 kernel 需要的资源调整block
• 保证 block 的数目远大于 SM 的数目
• 多做实验来挖掘出最好的配置

七、Synchronize

同步是并行编程的一个普遍的问题。在CUDA中，有多种方式实现同步：

1. System-level：等待所有 device 的工作完成
2. Block-level：等待 device 中 block 的所有 thread 执行到某个点
3. Stream-level：阻塞主机线程，直到指定流中的所有操作都完成

因为 CUDA API 和 host 代码是异步的，cudaDeviceSynchronize 可以用来停住 CUP 等待 CUDA 中的操作完成：cudaError_t cudaDeviceSynchronize(void); 

因为 block 中的 thread 执行顺序不定，CUDA 提供了一个 function 来同步 block 中的 thread：__device__ void __syncthreads(void); 当该函数被调用，block中的每个thread都会等待所有其他thread执行到某个点来实现同步

cudaStreamSyncronize(stream); 流同步,可导致主机代码阻塞,直到给定的流完成其操作为止