大规模并行处理器编程：实践导向方法

📚 内容概要

本课程全面介绍了使用CUDA C环境进行GPU计算与并行编程的基础知识。内容涵盖GPU架构、数据并行性、线程管理、内存优化以及高级性能考量，并通过磁共振成像（MRI）重建和分子可视化等实际案例加以说明。

掌握高性能并行计算的艺术，通过实践导向的指南学习CUDA与GPU架构。

作者： 大卫·B·基尔克（David B. Kirk），胡文梅（Wen-mei W. Hwu）

致谢： 伊恩·巴克（Ian Buck）、约翰·尼克尔斯（John Nickolls）、NVIDIA DevTech团队、黄仁勋（Jensen Huang）、戴维·吕布克（David Luebke）、比尔·比恩（Bill Bean）、西蒙·格林（Simon Green）、马克·哈里斯（Mark Harris）、曼朱·赫吉（Manju Hedge）、纳迪姆·莫汉默德（Nadeem Mohammad）、布伦特·奥斯特（Brent Oster）、彼得·肖利（Peter Shirley）、埃里克·杨（Eric Young）和西里尔·泽勒（Cyril Zeller）。

🎯 学习目标

1. 区分多核CPU与多核GPU在设计哲学和性能演进路径上的差异。
2. 识别现代GPU架构的关键组件，包括流式多处理器（SMs）和内存结构。
3. 应用阿姆达尔定律计算理论加速比，并识别顺序瓶颈的影响。
4. 对比固定功能流水线与可编程统一处理器阵列之间的架构差异。
5. 解释“GPGPU”作为过渡阶段的作用，以及早期着色器编程模型的局限性。
6. 分析原子操作、屏障同步和双精度支持等硬件特性如何推动向可扩展通用计算的转变。
7. 识别并利用矩阵-矩阵乘法算法中的数据并行性。
8. 实现设备内存管理，包括分配、主机与设备间的数据传输，以及释放。
9. 使用适当的线程索引和网格/块配置构建并启动CUDA内核。
10. 设计多维线程层次结构（网格与块），将复杂数据结构映射到GPU硬件上。

🔹 第1课：并行计算与GPU架构导论

概述： 本课探讨从串行计算向并行计算的根本转变，这一转变源于CPU与GPU设计哲学的分化。学生将研究“多核”与“众核”的发展轨迹，理解使GPU实现巨大吞吐量的硬件架构，并学习阿姆达尔定律对加速比的数学约束。

学习成果：

• 区分多核CPU与多核GPU在设计哲学和性能演进路径上的差异。
• 识别现代GPU架构的关键组件，包括流式多处理器（SMs）和内存结构。
• 应用阿姆达尔定律计算理论加速比，并识别顺序瓶颈的影响。

🔹 第2课：GPU计算的演变与未来展望

概述： 本课追溯图形处理单元（GPU）的架构发展历程，从最初用于三角形渲染的专用固定功能硬件，演变为如今强大、统一且通用的并行处理器。学生将探索从刚性图形流水线到可编程着色器的转变，了解GPGPU运动的兴起，以及驱动当前科学与工程模拟的现代可扩展架构。

学习成果：

• 对比固定功能流水线与可编程统一处理器阵列之间的架构差异。
• 解释“GPGPU”作为过渡阶段的作用，以及早期着色器编程模型的限制。
• 分析原子操作、屏障同步和双精度支持等硬件特性如何促成向可扩展通用计算的转变。

🔹 第3课：CUDA程序结构与内存管理

概述： 本课介绍CUDA程序的基本架构，强调主机（CPU）与设备（GPU）执行之间的区别。学生将学会识别矩阵运算中的数据并行性，使用CUDA API管理独立的内存空间，并通过单程序多数据（SPMD）风格组织并行执行，形成网格、块与线程的层级结构。

学习成果：

• 识别并利用矩阵-矩阵乘法算法中的数据并行性。
• 实现设备内存管理，包括分配、主机与设备间的数据传输，以及释放。
• 使用适当的线程索引和网格/块配置构建并启动CUDA内核。

🔹 第4课：高级CUDA线程与调度

概述： 本课深入探讨CUDA中线程的层级组织，重点分析多维索引如何映射到物理数据与硬件资源。内容详述屏障同步机制与透明可扩展性，最终总结出用于实现高性能计算中容忍延迟的线程分配与基于瓦片（warp）调度的架构原则。

学习成果：

• 设计多维线程层次结构（网格与块），将复杂数据结构映射到GPU硬件。
• 实现使用内置CUDA变量（blockIdx、threadIdx、blockDim）进行精确数据索引。
• 应用屏障同步，确保数据完整性的同时，在不同GPU架构间保持透明可扩展性。

🔹 第5课：内存优化与共享内存分块

概述： 本课探讨内存带宽与资源限制在并行计算中作为主要瓶颈的作用。详细说明“分块”技术如何减少全局内存流量，并解释同步屏障（__syncthreads()）的关键作用，以及在寄存器与共享内存之间做出战略性选择以优化性能。

学习成果：

• 分析寄存器与共享内存容量如何决定内核中的并行度（占用率）。
• 量化通过分块技术实现的全局内存带宽消耗减少量。
• 识别同步函数在共享内存访问期间维护数据完整性的必要性。

🔹 第6课：性能分析与SIMT执行

概述： 本课探讨优化CUDA内核所必需的架构与算法考量。从基本执行模型——特别是单指令多线程（SIMT）单元与瓦片划分——过渡到高级性能调优技术，包括内存合并、分块矩阵乘法，以及流式多处理器（SM）资源的动态分区。

学习成果：

• 分析多维线程块如何映射到硬件的线性瓦片执行顺序。
• 评估并最小化并行归约算法中的控制流分歧。
• 通过实现内存合并与分块数据访问模式，优化全局内存带宽。

🔹 第7课：浮点数算术与数值精度

概述： 本课介绍浮点数的基本架构，聚焦于IEEE 754标准的组成部分：符号位、偏移编码指数和规格化尾数。学生将探索这些位模式如何映射到离散数轴，并分析这种表示方式的局限性对大规模求和等复杂算法精度的影响。

学习成果：

• 解析浮点格式，通过规格化表示与偏移编码从位模式计算数值。
• 可视化可表示数在数轴上的分布，并解释指数与尾数位分配对结果的影响。
• 使用单位最后一位（ULP）量化数值不准确性，并识别不同舍入模式对误差的贡献。

🔹 第8课：案例研究——并行化MRI重建

概述： 本课探讨在GPU上并行化先进的磁共振成像（MRI）重建过程。重点针对非笛卡尔轨迹的迭代重建，具体通过循环变换、常量内存管理、数据布局重组，以及使用硬件加速的三角函数来优化计算密集型的 F^H d 内核。

学习成果：

• 理解从基于傅里叶变换的笛卡尔重建转向适用于非笛卡尔k空间数据的迭代线性求解算法的过程。
• 应用循环分裂与循环交换，将顺序C代码转换为适合大规模CUDA线程映射的结构。
• 通过常量内存分块和结构体数组（AoS）数据布局优化内存吞吐量。

🔹 第9课：案例研究——分子可视化与多GPU执行

概述： 本课探讨GPU计算在分子可视化中的实际应用，特别是使用直接库仑求和（DCS）方法计算静电势图。学生将从基础内核实现逐步过渡到高度优化的版本，利用指令展开、内存合并和填充技术提升性能。

学习成果：

• 使用CUDA常量内存与全局内存延迟隐藏技术实现直接库仑求和（DCS）内核。
• 通过指令展开和重用常见坐标计算优化内核性能。
• 应用内存合并与填充策略，对齐GPU全局内存访问以实现最大带宽。

🔹 第10课：计算思维与并行算法选择

概述： 本课探讨从串行思维向并行问题解决的转变，重点关注并行编程的目标以及算法的战略选择。学生将学习如何将问题分解为可并行化的单元，运用计算思维弥合领域科学与硬件架构之间的差距，并评估算法性能。

学习成果：

• 识别并行编程的主要目标，并使用阿姆达尔定律计算理论加速比。
• 区分任务级与数据级分解，应用以原子为中心（散射）与以网格为中心（聚集）的策略。
• 根据内存带宽、计算复杂度和架构约束等标准，评估并选择并行算法。

🔹 第11课：OpenCL编程模型入门

概述： 本课介绍OpenCL作为一种异构并行计算框架，重点在于其数据并行模型与分层硬件抽象。学生将学习如何将OpenCL的NDRange与内存结构映射到CUDA对应架构，并掌握通过动态编译模型管理主机端设备的方法。

学习成果：

• 将OpenCL的并行性与内存层次结构映射到特定于CUDA的架构（例如将工作组映射到块，本地内存映射到共享内存）。
• 实现OpenCL内核函数，并使用上下文与命令队列管理主机端执行环境。
• 执行动态编译流程，实现在运行时从源代码构建内核。

🔹 第12课：现代GPU特性与未来展望

概述： 本课探讨GPU在架构与功能上的演进，重点关注向复杂内存管理、增强内核执行能力及核心性能提升的转变。学生将分析统一设备内存空间与内核级函数调用等功能如何使GPU转变为通用处理器。

学习成果：

• 解释内存架构演进的意义，以及向64位统一设备内存空间迈进的重要性。
• 分析增强型原子操作与内核级函数调用如何支持复杂数据结构与算法的实现。
• 评估现代GPU环境中同时执行多个内核、双精度性能提升以及控制流效率带来的性能影响。