# 案例与习题 作者：Jason D. Bakos。 ## 案例：低算术强度机器学习任务中的 GPU 与向量处理器 本案例涉及的概念： - 算术强度与性能瓶颈 - GPU（SIMT）编程模型与并行性 - 向量处理器编程模型与并行性 考虑一个执行神经网络批量归一化（batch normalization）层和整流线性单元（ReLU）层的内核。把这两层合在一起，对应的 C 代码如下。 ~~~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 {#exercise-4-1} **[5]〈4.3〉** 从输入和输出数组每访问一个字节，需要执行多少次算术运算？ ### 习题 4.2 {#exercise-4-2} **[15]〈4.4〉** 用 PTX 实现这个循环体。假设 `in` 和 `out` 数组存放在全局内存中，`mean`、`stddev`、`scale` 和 `offset` 数组存放在共享内存中。 ### 习题 4.3 {#exercise-4-3} **[5]〈4.4〉** 假设在你的 GPU 上，加载指令的延迟为 10 个周期，其他所有浮点指令的延迟均为 4 个周期。GPU 核心上必须执行多少个 Warp，才能隐藏所有指令延迟，也就是让核心始终保持 100% 利用率？假设核心的寄存器数量无限，存储器系统的带宽也无限。若每个 Warp 有 32 个线程，对应的线程总数是多少？ ### 习题 4.4 {#exercise-4-4} **[4]〈4.4〉** 假设代码结构如下，并且在一个 Warp 内，\(P(\mathrm{cond1})=35\%\)、\(P(\mathrm{cond2})=10\%\)、\(P(\mathrm{cond3})=40\%\)。`codeblock1` 的延迟为 100 个周期，`codeblock2` 为 200 个周期，`codeblock3` 为 300 个周期。预期执行时间是多少？ ~~~c if (cond1) { if (cond2) { codeblock1 } else { codeblock2 } } else { if (cond3) { codeblock3 } else { codeblock4 } } ~~~ > **译注：** 原题代码含 `codeblock4`，但题干没有给出其延迟，也没有说明各线程条件之间的相关性。 ### 习题 4.5 {#exercise-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 指令： ~~~asm vfmax.vf vd, vs2, rs1, vm # 向量—标量最大值，执行 vd(*) = max([vs2(*)], [rs1]) # 将 vm 设为 1 可禁用掩码 ~~~ ### 习题 4.6 {#exercise-4-6} **[10]〈4.2〉** 假设有两个加载／存储单元和五个向量浮点单元。重新安排 e 部分的代码，把指令组织成编队。 ### 习题 4.7 {#exercise-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 {#exercise-4-8} **[10/20/20/15/15]〈4.2〉** 考虑下面的代码，它将两个包含单精度复数的向量相乘： ~~~c 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 {#exercise-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 {#exercise-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 {#exercise-4-11} **[15/25/25]〈4.4、4.5〉** 4.5 节讨论了归约操作，即通过反复应用某个运算，把一个向量缩减为一个标量。归约是循环递推的一种特殊形式。例如： ~~~c dot=0.0; for (i=0; i < 64; i++) dot = dot + a[i] * b[i]; ~~~ 向量化编译器可以应用一种名为**标量扩展（scalar expansion）**的变换，把 `dot` 扩展为向量，并拆分循环，使乘法能够用向量运算完成，而归约保留为单独的标量运算： ~~~c 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 节所述，如果允许浮点加法满足结合律，就可以用多种技术并行执行归约。 ~~~c 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 {#exercise-4-12} **[10/10/10/10]〈4.3〉** 下面的内核实现了时域有限差分（finite-difference time-domain，FDTD）方法的一部分，用于计算三维空间中的 Maxwell 方程。这是 SPEC06fp 某个基准的一部分： ~~~c 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 {#exercise-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 {#exercise-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 {#exercise-4-15} **[10/15/15]〈4.5〉** 本题考察几个循环，并分析其并行化潜力。 a. **[10]〈4.5〉** 下面的循环有循环携带依赖吗？ ~~~c for (i=0; i < 100; i++) { A[i] = B[2*i+4]; B[4*i+5] = A[i]; } ~~~ b. **[15]〈4.5〉** 找出下面循环中的所有真依赖、输出依赖和反依赖。通过重命名消除输出依赖和反依赖。 ~~~c 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〉** 考虑下面的循环： ~~~c 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 {#exercise-4-16} **[10]〈4.4〉** 列出并描述至少四个影响 GPU 内核性能的因素。换句话说，由内核代码引起的哪些运行时行为，会使内核执行过程中的资源利用率下降？ ### 习题 4.17 {#exercise-4-17} **[10]〈4.4〉** 假设有一个 GPU，具有如下特征： - 时钟频率为 1.5 GHz。 - 包含 16 个 SIMD 处理器，每个处理器包含 16 个单精度浮点单元。 - 片外存储器带宽为 100 GB/s。 假设所有存储器延迟都能被隐藏，不考虑存储器带宽时，该 GPU 的单精度浮点峰值吞吐量是多少 GFLOP/s？考虑存储器带宽限制后，能否持续保持这一吞吐量？ ### 习题 4.18 {#exercise-4-18} **[60]〈4.4〉** 在这个编程练习中，你将编写一个包含大量数据级并行性、同时也包含条件执行行为的 CUDA 内核，并分析其行为。使用 NVIDIA CUDA Toolkit，以及不列颠哥伦比亚大学的 GPU-SIM（原书链接：[gpgpu-sim.org](http://www.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〉** 实现内核的改进版本，使片外存储器访问能够合并，并观察运行性能的变化。