cuda4

性能

Warps and SIMD

一个线程 block 由 32 个线程 warp 组成
扭曲在多处理器上以物理方式并行执行 (SIMD)

Warp 是调度单位

---

控制分支

• 当 warp 中的线程通过做出不同的控制决策而采取不同的控制流路径时，就会发生控制分支
  • 一些采用 then 路径，另一些采用 if 语句的 else 路径
  • 一些线程与其他线程采用不同数量的循环迭代
• 采取不同路径的线程的执行在当前的 GPU 中被序列化
  • 一个 warp 中的线程所采用的控制路径一次遍历一个，直到不再存在。
  • 考虑嵌套控制流语句时，不同路径的数量可能很大

控制分支例子

当分支或循环条件是线程索引的函数时，可能会出现分歧

具有分歧的内核语句示例：
如果 (threadIdx.x > 2) { }
这为块中的线程创建了两个不同的控制路径
决策粒度 < 扭曲大小； 线程 0、1 和 2 遵循与第一个 warp 中的其余线程不同的路径

没有发散的例子：
如果 (blockIdx.x > 2) { }
决策粒度是块大小的倍数； 任何给定 warp 中的所有线程都遵循相同的路径

---

控制分支的性能影响

• 边界条件检查对于并行代码的完整功能和健壮性至关重要
  • 分块矩阵乘法内核有很多边界条件检查
  • 令人担忧的是，这些检查可能会导致性能显着下降

if(Row < Width && p * TILE_WIDTH+tx < Width) {
   ds_M[ty][tx] = M[Row * Width + p * TILE_WIDTH + tx];
} else {
	 ds_M[ty][tx] = 0.0;
}

if (p*TILE_WIDTH+ty < Width && Col < Width) {
   ds_N[ty][tx] = N[(p*TILE_WIDTH + ty) * Width + Col];
} else {
   ds_N[ty][tx] = 0.0;

加载 M Tiles 时的两种块

1. 直到最后一个阶段，其 tiles 都在有效范围内的块。
2. 方块有部分一直在有效范围之外

控制分支影响分析

假设 16x16 tiles 和线程块
每个线程块有 8 个 warp (256/32)
假设 100x100 的方阵
每个线程将经历 7 个阶段（上限为 100/16）

有 49 个线程块（每个维度 7 个）

加载 M tiles 的控制分支

TYPE1

假设 16x16 TILES 和线程块
每个线程块有 8 个 WARP (256/32)
假设 100x100 的方阵
每个经线将经历 7 个阶段（100/16 的上限）

有 42（$67$）个类型 1 块，总共有 336（$842$）条 warps
它们都有 7 个阶段，因此有 2,352 (336*7) 个 WARP 阶段
经线只有在最后阶段才有控制发散
336 个经线阶段有控制分支

7 个阶段：每行取七次，最后一次不完整

只考虑 Warp 不考虑 Block 不完整：因为 Block 不完整会导致整个 Warps 都不取，也就不存在分支

336 个阶段：6*7*8*1 6*7 个 block，每个 8 个 Warp

TYPE2

类型 2：分配加载底部 TILES 的 7 个块，共 56（$87$）个扭曲
它们都有 7 个阶段，所以有 392 ($567$) 个 WARP 阶段
每个类型 2 块中的前 2 个 WARP 将保持在有效范围内，直到最后一个阶段
剩余的 6 个 WARP 不在有效范围内
所以，只有 14 (2*7) 个经线阶段有控制分支

14 个阶段：2*7*1 7 个 block，每个 2 个 Warp

2 个 Warp：两个横排，一个横排 16 个

在大矩阵情况下，对于性能影响很小

控制分支总体影响

类型 1： 块：2,352 个 warp 阶段中的 336 个具有控制分支
类型 2： 块：392 个 warp 阶段中有 14 个具有控制分支
性能影响预计小于 12% (350/2,944 或 (336+14)/(2352+14))

Add。

加载 N 个 TILEs 时控制发散的影响计算有些不同，留作练习

估计的性能影响取决于数据。
对于较大的矩阵，影响将显着较小

一般来说，控制发散对大型输入数据集的边界条件检查的影响应该是微不足道的
应该毫不犹豫地使用边界检查来确保完整的功能

内核中充满控制流结构的事实并不意味着会出现严重的控制发散

我们将在 Parallel Algorithm Patterns 模块中介绍一些自然会导致控制发散（例如并行缩减）的算法模式

---

并行规约

划分和总结

将数据集分成更小的块
让每个线程处理一个块
使用归约树将每个块的结果汇总为最终答案

将大的问题分解成小的问题，让每个线程负责一个问题，并利用一棵树将结果归约为最终结果。

Reduction Conputation

规约将一组输入的数组汇总成一个值，例如：

• 求最值
• 求和
• 积

算法复杂度 o(N)

并行求和规约

每个线程负责两个值的求和，需要 n/2 个线程，执行 log(n)次。

in-place 不使用辅助变量来转换输入数据结构

一个简单的数据映射线程

每个线程负责部分和向量的偶数索引位置（位置责任）
每一步后，不再需要一半的线程
输入之一总是来自责任地点
在每一步中，其中一个输入来自越来越远的距离

__shared__ float partialSum[2*BLOCK_SIZE];

unsigned int t = threadIdx.x;
unsigned int start = 2*blockIdx.x*blockDim.x;
partialSum[t] = input[start + t];
partialSum[blockDim.x+t] = input[start + blockDim.x+t];

for (unsigned int stride = 1; stride <= blockDim.x;  stride *= 2)
{
  __syncthreads();
  if (t % stride == 0)
    partialSum[2*t]+= partialSum[2*t+stride];
}

同步是因为需要在进行下一步前，获得上一步的所有结果，下一步的操作数来源是新的

求和完成后，如果 Block 非常多，宿主代码可以迭代启动另一个内核进行求和；若较少，则可以传回主机 加和，或利用原子操作累加到全局变量中。

优化

每次迭代后 Warp 中真正参与运算的线程很少，资源利用率非常低，在 5 次之后每个 Warp 中只有一个线程在运行但却占用了整个 Warp 的资源 通过改变索引改善，使得部分和压缩在数组的前面位置

在一些算法中，可以改变索引的使用来改善发散行为
交换和结合运算符
始终将部分和压缩到 partialSum[] 数组中的前面位置
保持活动线程连续

更好的核函数

for (unsigned int stride = blockDim.x;
	  stride > 0;  stride /= 2)
{
  __syncthreads();
  if (t < stride)
	  partialSum[t] += partialSum[t+stride];
}

---

内存并行

全局内存（DRAM）带宽

DRAM 核心阵列组织

DRAM 核心阵列很慢

–DDR: Core speed = ½ interface speed

–DDR2/GDDR3: Core speed = ¼ interface speed

–DDR3/GDDR4: Core speed = ⅛ interface speed

DRAM Bursting （突发）

通过将 N 倍位宽的数据加载至缓冲区，随后以 N 步读出（仅适用于连续地址）

复数 Bank 时类似

将内存地址划分为几个不同的区域，当一个地址被读取，整个区域被送出。

内存合并

因此，当一个 Warp 中的所有线程执行一个 load 时，且访问位在同一个突发区域中时，只会发出一个读取 指令，且访问合并。快。若不是这样，就会发出多个请求，并且一些读出的数据被丢弃。

如果数组访问中的索引采用以下形式，则扭曲中的访问是对连续位置的访问

• _A[(expression with terms independent of threadIdx.x) + threadIdx.x];_（ 中间英文：具有独立项的表达式）