案例与习题

作者：Jason D. Bakos。

案例：低算术强度机器学习任务中的 GPU 与向量处理器

本案例涉及的概念：

• 算术强度与性能瓶颈
• GPU（SIMT）编程模型与并行性
• 向量处理器编程模型与并行性

考虑一个执行神经网络批量归一化（batch normalization）层和整流线性单元（ReLU）层的内核。把这两层合在一起，对应的 C 代码如下。

for (uint32_t i=0; i < n_channels; i++) {
    temp1 = max(in[i],0);
    temp2 = (temp1 - mean[i]) / stddev[i];
    out[i] = temp2 * scale[i] + offset[i];
}

假设所有变量均为 float16 类型。

习题 4.1

[5]〈4.3〉 从输入和输出数组每访问一个字节，需要执行多少次算术运算？

习题 4.2

[15]〈4.4〉 用 PTX 实现这个循环体。假设 in 和 out 数组存放在全局内存中，mean、stddev、scale 和 offset 数组存放在共享内存中。

习题 4.3

[5]〈4.4〉 假设在你的 GPU 上，加载指令的延迟为 10 个周期，其他所有浮点指令的延迟均为 4 个周期。GPU 核心上必须执行多少个 Warp，才能隐藏所有指令延迟，也就是让核心始终保持 100% 利用率？假设核心的寄存器数量无限，存储器系统的带宽也无限。若每个 Warp 有 32 个线程，对应的线程总数是多少？

习题 4.4

[4]〈4.4〉 假设代码结构如下，并且在一个 Warp 内，$P(\mathrm{cond1})=35\%$、$P(\mathrm{cond2})=10\%$、$P(\mathrm{cond3})=40\%$。codeblock1 的延迟为 100 个周期，codeblock2 为 200 个周期，codeblock3 为 300 个周期。预期执行时间是多少？

if (cond1) {
    if (cond2) {
        codeblock1
    } else {
        codeblock2
    }
} else {
    if (cond3) {
        codeblock3
    } else {
        codeblock4
    }
}

译注： 原题代码含 codeblock4，但题干没有给出其延迟，也没有说明各线程条件之间的相关性。

习题 4.5

[15]〈4.2〉 用 RV64V 汇编实现 a 部分的代码。假设该体系结构支持 16 位浮点数，寄存器初值如下：

• x1：in 数组的基地址
• x2：mean 数组的基地址
• x3：stddev 数组的基地址
• x4：scale 数组的基地址
• x5：offset 数组的基地址
• x6：out 数组的基地址
• x7：n_channels

你可能需要下面这条未列入图 4.2 的 RV64V 指令：

vfmax.vf vd, vs2, rs1, vm
# 向量—标量最大值，执行 vd(*) = max([vs2(*)], [rs1])
# 将 vm 设为 1 可禁用掩码

习题 4.6

[10]〈4.2〉 假设有两个加载／存储单元和五个向量浮点单元。重新安排 e 部分的代码，把指令组织成编队。

习题 4.7

[5]〈4.2〉 向量长度为 32 时，f 部分的向量指令序列需要多少个 chime？每 FLOP 的周期数是多少？

译注： 原题 4.5、4.6、4.7 分别引用“a 部分”“e 部分”“f 部分”，但本版将前面的题目排为独立编号 4.1–4.7，没有这些字母小问。结合连续题干，它们分别承接案例代码、习题 4.5 与习题 4.6。

习题 4.8

[10/20/20/15/15]〈4.2〉 考虑下面的代码，它将两个包含单精度复数的向量相乘：

for (i=0; i < 300; i++) {
    c_re[i] = a_re[i] * b_re[i] - a_im[i] * b_im[i];
    c_im[i] = a_re[i] * b_im[i] + a_im[i] * b_re[i];
}

习题 4.9

假设处理器工作频率为 700 MHz，最大向量长度为 64。加载／存储单元的启动开销为 15 个周期，乘法单元为 8 个周期，加／减法单元为 5 个周期。

a. [10]〈4.3〉 这个内核的算术强度是多少？说明理由。

b. [20]〈4.2〉 使用 strip mining，将这个循环转换为 RV64V 汇编代码。

c. [20]〈4.2〉 假设支持链接，并且只有一条存储器流水线，需要多少个 chime？包括启动开销在内，每个复数结果需要多少个时钟周期？

d. [15]〈4.2〉 如果向量指令序列采用链接，包括开销在内，每个复数结果需要多少个时钟周期？

e. [15]〈4.2〉 现在假设处理器有三条存储器流水线，并支持链接。如果循环中的访问没有存储体冲突，每个结果需要多少个时钟周期？

译注： 原页将复数乘法代码标为习题 4.8，将后续假设和 a–e 小问标为 4.9。c、d 两问原文都写了链接条件，条件存在重复，未擅自把其中一问改成“不支持链接”。

习题 4.10

[30]〈4.2、4.3、4.4〉 本题比较向量处理器，与一个由标量处理器和 GPU 协处理器组成的混合系统的性能。在混合系统中，主机处理器的标量性能优于 GPU，所以所有标量代码均在主机处理器上执行，所有向量代码均在 GPU 上执行。我们把第一个系统称为向量计算机，第二个系统称为混合计算机。

假设目标应用包含一个向量内核，其算术强度为每访问一个 DRAM 字节执行 0.5 FLOP。不过，应用还有一个标量部分，必须在内核执行之前和之后分别准备输入向量与输出向量。对于一个样本数据集，标量代码部分在向量处理器，以及混合系统的主机处理器上，都需要 400 ms 执行时间。内核读取的输入向量共包含 200 MB 数据，输出数据为 100 MB。向量处理器的峰值存储器带宽为 30 GB/s，GPU 的峰值存储器带宽为 150 GB/s。混合系统还有额外开销：在调用内核前后，需要在主机内存与 GPU 本地内存之间传送所有输入向量。混合系统的直接存储器访问（DMA）带宽为 10 GB/s，平均延迟为 10 ms。假设向量处理器和 GPU 的性能都受存储器带宽限制。计算这两台计算机执行该应用所需的时间。

译注： 原题在“内核调用前后”都写“所有输入向量”，并未在此处区分输入传入与输出传回；这里忠实保留。此处 GPU local memory 指与主机内存相对的 GPU 本地内存，不等同于前文 NVIDIA 为每线程私有地址空间使用的 Local Memory 名称。

习题 4.11

[15/25/25]〈4.4、4.5〉 4.5 节讨论了归约操作，即通过反复应用某个运算，把一个向量缩减为一个标量。归约是循环递推的一种特殊形式。例如：

dot=0.0;
for (i=0; i < 64; i++) dot = dot + a[i] * b[i];

向量化编译器可以应用一种名为**标量扩展（scalar expansion）**的变换，把 dot 扩展为向量，并拆分循环，使乘法能够用向量运算完成，而归约保留为单独的标量运算：

for (i=0; i < 64; i++) dot[i] = a[i] * b[i];
for (i=1; i < 64; i++) dot[0] = dot[0] + dot[i];

如 4.5 节所述，如果允许浮点加法满足结合律，就可以用多种技术并行执行归约。

unsigned int tid = threadIdx.x;
for(unsigned int s=1; s < blockDim.x; s * = 2) {
    if ((tid % (2*s)) == 0) {
        sdata[tid] += sdata[tid + s];
    }
    __syncthreads();
}

重写这个循环，使其符合这些准则，并消除取模运算符。假设每个 Warp 有 32 个线程；只要同一 Warp 中两个或更多线程所访问下标除以 32 的余数相同，就会发生存储体冲突。

a. [15]〈4.4、4.5〉 一种方法称为递推倍增（recurrence doubling），它依次对越来越短的向量进行相加，例如先把两个 32 元素向量相加，再把两个 16 元素向量相加，依此类推。写出以这种方式执行第二个循环的 C 代码。

b. [25]〈4.4、4.5〉 在某些向量处理器中，向量寄存器内的单个元素可以寻址。因此，一个向量运算的两个操作数可以来自同一向量寄存器的不同部分。利用这种灵活性，还可以采用另一种名为**部分和（partial sums）**的归约方案：将向量归约为 $m$ 个和，其中 $m$ 是向量功能单元的总延迟，包括操作数读写的时间。

假设 RV64 向量寄存器可以按元素寻址，例如，可以用操作数 V1(16) 启动向量运算，表示输入操作数从第 16 个元素开始。另假设加法的总延迟为 8 个周期，包括读取操作数和写入结果的时间。写出一个 RV64V 代码序列，将 V1 的内容归约为 8 个部分和。

c. [25]〈4.4、4.5〉 在 GPU 上执行归约时，为输入向量的每个元素关联一个线程。首先，每个线程把相应的值写入共享内存。然后，每个线程进入一个循环，将输入值两两相加。每次迭代后，元素数量减半，活跃线程数也随之减半。为了尽可能提高归约性能，整个循环过程中应尽量增加所有线程都活跃的 Warp 数量。换句话说，活跃线程应当连续。此外，每个线程对共享数组的索引方式还应避免共享内存中的存储体冲突。下面的循环只违反第一条准则，并使用了对 GPU 而言代价很高的取模运算符。

译注： 原书把 c 小问所称的“下面的循环”，以及要求改写该循环的段落，排在 a 小问之前。原页循环更新式写作 s * = 2，其中空格拆开了复合赋值运算符；可编译的 C/CUDA 写法应为 s *= 2。

习题 4.12

[10/10/10/10]〈4.3〉 下面的内核实现了时域有限差分（finite-difference time-domain，FDTD）方法的一部分，用于计算三维空间中的 Maxwell 方程。这是 SPEC06fp 某个基准的一部分：

for (int x=0; x < NX - 1; x++) {
    for (int y=0; y < NY - 1; y++) {
        for (int z=0; z < NZ - 1; z++) {
            int index = x*NY*NZ + y*NZ + z;
            if (y > 0 && x > 0) {
                material = IDx[index];
                dH1 = (Hz[index] - Hz[index-incrementY])/dy[y];
                dH2 = (Hy[index] - Hy[index-incrementZ])/dz[z];
                Ex[index] = Ca[material]*Ex[index]+Cb[material]* (dH2 - dH1);
}}}}

假设 dH1、dH2、Hy、Hz、dy、dz、Ca、Cb 和 Ex 都是单精度浮点数组，IDx 是 unsigned int 数组。

a. [10]〈4.3〉 这个内核的算术强度是多少？

b. [10]〈4.3〉 这个内核适合向量执行或 SIMD 执行吗？为什么？

c. [10]〈4.3〉 假设在存储器带宽为 30 GB/s 的处理器上执行这个内核。它受存储器限制，还是受计算能力限制？

d. [10]〈4.3〉 假设该处理器的峰值计算吞吐量为 85 GFLOP/s，建立它的 Roofline 模型。

译注： 原题把 dH1、dH2 列为数组，但代码中按没有下标的标量使用。

习题 4.13

[10/15]〈4.4〉 假设某 GPU 体系结构包含 10 个 SIMD 处理器。每条 SIMD 指令的宽度为 32，每个 SIMD 处理器包含 8 条用于单精度算术与加载／存储指令的通道，即每条没有发生分歧的 SIMD 指令每 4 个周期能够产生 32 个结果。假设某内核中的分歧分支使平均只有 80% 的线程处于活跃状态。执行的全部 SIMD 指令中，70% 是单精度算术指令，20% 是加载／存储指令。由于部分存储器延迟未能被隐藏，假设平均 SIMD 指令发射率为 0.85。GPU 的时钟频率为 1.5 GHz。

习题 4.14

下面各项改进带来的吞吐量加速比分别是多少？

a. [10]〈4.4〉 计算上述内核在该 GPU 上的吞吐量，单位为 GFLOP/s。

b. [15]〈4.4〉 假设有以下选择：

1. 把单精度通道数增加到 16。
2. 把 SIMD 处理器数增加到 15。假设这项改动不影响任何其他性能指标，且代码能够扩展到新增的处理器。
3. 增加一个缓存，将存储器延迟有效降低 40%，使指令发射率提高到 0.95。

译注： 原页将体系结构假设排为习题 4.13，将问题和 a、b 小问排为 4.14；此处沿用原编号，不把它们合并或重新编号。

习题 4.15

[10/15/15]〈4.5〉 本题考察几个循环，并分析其并行化潜力。

a. [10]〈4.5〉 下面的循环有循环携带依赖吗？

for (i=0; i < 100; i++) {
    A[i] = B[2*i+4];
    B[4*i+5] = A[i];
}

b. [15]〈4.5〉 找出下面循环中的所有真依赖、输出依赖和反依赖。通过重命名消除输出依赖和反依赖。

for (i=0; i < 100; i++) {
    A[i] = A[i] * B[i]; /* S1 */
    B[i] = A[i] + c;    /* S2 */
    A[i] = C[i] * c;    /* S3 */
    C[i] = D[i] * A[i]; /* S4 */
}

c. [15]〈4.5〉 考虑下面的循环：

for (i=0; i < 100; i++) {
    A[i] = A[i] + B[i];   /* S1 */
    B[i+1] = C[i] + D[i]; /* S2 */
}

S1 与 S2 之间有依赖吗？这个循环是并行的吗？如果不是，说明如何使其并行。

习题 4.16

[10]〈4.4〉 列出并描述至少四个影响 GPU 内核性能的因素。换句话说，由内核代码引起的哪些运行时行为，会使内核执行过程中的资源利用率下降？

习题 4.17

[10]〈4.4〉 假设有一个 GPU，具有如下特征：

• 时钟频率为 1.5 GHz。
• 包含 16 个 SIMD 处理器，每个处理器包含 16 个单精度浮点单元。
• 片外存储器带宽为 100 GB/s。

假设所有存储器延迟都能被隐藏，不考虑存储器带宽时，该 GPU 的单精度浮点峰值吞吐量是多少 GFLOP/s？考虑存储器带宽限制后，能否持续保持这一吞吐量？

习题 4.18

[60]〈4.4〉 在这个编程练习中，你将编写一个包含大量数据级并行性、同时也包含条件执行行为的 CUDA 内核，并分析其行为。使用 NVIDIA CUDA Toolkit，以及不列颠哥伦比亚大学的 GPU-SIM（原书链接：gpgpu-sim.org）或 CUDA Profiler，编写并编译一个 CUDA 内核：在 $256\times256$ 的棋盘上执行 Conway 生命游戏的 100 次迭代，并将棋盘的最终状态返回主机。假设棋盘由主机初始化。为每个细胞分配一个线程。确保每次游戏迭代后都加入屏障。使用以下游戏规则：

• 任何活细胞，如果活邻居少于两个，就死亡。
• 任何活细胞，如果有两个或三个活邻居，就存活到下一代。
• 任何活细胞，如果活邻居多于三个，就死亡。
• 任何死细胞，如果恰好有三个活邻居，就变成活细胞。

完成内核后，回答以下问题：

a. [60]〈4.4〉 使用 -ptx 选项编译代码，检查内核的 PTX 表示。该内核的 PTX 实现包含多少条 PTX 指令？内核中的条件执行部分包含分支指令，还是仅包含谓词化的非分支指令？

b. [60]〈4.4〉 在模拟器中执行代码后，动态指令数是多少？达到的每周期指令数（IPC）或指令发射率是多少？动态指令中，控制指令、算术逻辑单元（ALU）指令和存储器指令分别占多少？是否存在共享内存存储体冲突？有效片外存储器带宽是多少？

c. [60]〈4.4〉 实现内核的改进版本，使片外存储器访问能够合并，并观察运行性能的变化。