第2章 GPU架构 - SIMT的精髓 (GPU Architecture - The Essence of SIMT)

一、导读

2.1 本章学习目标

通过本章的学习，读者将能够：

深入理解SIMT（Single Instruction, Multiple Threads）执行模型的工作原理和硬件实现

掌握GPU微架构的核心组件：SM（Streaming Multiprocessor）、warp调度器、执行单元

理解warp的概念及其在GPU并行执行中的关键作用

分析分支分歧对SIMT执行效率的影响机制

掌握GPU指令流水线的特点和与CPU流水线的差异

理解寄存器文件、共享内存、缓存子系统在GPU架构中的角色

能够根据GPU架构特点优化kernel代码的执行效率

了解NVIDIA不同代GPU架构的演进历程和技术创新

2.2 为什么本章如此重要

GPU架构是理解CUDA编程性能优化的基础。与CPU不同，GPU采用了截然不同的设计理念：牺牲单线程性能和控制灵活性，换取大规模并行计算能力。这种设计取舍体现在GPU的每一个架构细节中，从warp级别的线程组织到SM的资源分配，从简化的控制逻辑到高带宽的内存子系统。

对于CUDA程序员而言，理解GPU架构意味着能够预测代码的执行行为，识别性能瓶颈，并做出有针对性的优化决策。例如，知道warp的大小是32个线程，就能理解为什么分支分歧会导致性能下降；知道SM的寄存器文件有限，就能合理配置每个block的线程数量以最大化occupancy；知道共享内存的速度远快于全局内存，就能设计数据复用策略减少全局内存访问。

SIMT执行模型是GPU架构的灵魂。它结合了SIMD的高效率和多线程的灵活性，为程序员提供了独特的并行抽象。然而，SIMT也带来了新的挑战：warp内的线程必须协同工作，任何不协调的行为（如分支分歧、非合并内存访问）都会导致性能损失。只有深入理解SIMT的工作机制，才能编写出高效的GPU代码。

此外，GPU架构在不断演进。从早期的Tesla架构到最新的Hopper架构，每一代都引入了新的特性和优化。了解架构的演进历程有助于把握技术发展的脉络，理解新特性的设计动机，并为未来的编程挑战做好准备。

2.3 核心问题引导

什么是SIMT执行模型？它与SIMD和MIMD有何异同？

Warp是什么？为什么GPU选择32个线程作为一个warp？

SM（Streaming Multiprocessor）包含哪些核心组件？它们如何协同工作？

warp调度器如何选择下一个要执行的warp？零开销切换是如何实现的？

分支分歧如何影响warp的执行效率？硬件如何处理分支分歧？

GPU的寄存器文件与CPU有何不同？寄存器压力如何影响occupancy？

共享内存的物理实现是什么？bank conflict是如何产生的，如何避免？

NVIDIA的Volta、Ampere、Hopper架构各有哪些关键创新？

二、核心概念详解

2.1 SIMT执行模型的深度解析

SIMT（Single Instruction, Multiple Threads）是NVIDIA在CUDA架构中引入的执行模型，它是GPU并行计算的核心抽象。要深入理解SIMT，我们需要从多个维度进行分析。

基本定义：在SIMT模型中，多个线程组成一个执行组（在NVIDIA GPU中称为warp），warp内的所有线程在同一时钟周期执行相同的指令，但每个线程可以操作不同的数据，并且拥有独立的程序计数器（PC）和寄存器状态。

这与传统的SIMD（Single Instruction, Multiple Data）有本质区别。在纯SIMD模型中（如SSE、AVX），没有独立的线程概念，所有的数据lane共享同一个程序计数器，无法独立进行分支。当遇到条件语句时，SIMD通常使用掩码（mask/predicate）来选择性地更新某些lane的结果，或者串行执行各个分支路径。

SIMT的优势在于它为程序员提供了类似多线程的编程体验，同时保持了接近SIMD的执行效率。每个线程有自己的PC，可以独立地进行分支判断，这使得程序员可以用熟悉的if-else语句来表达控制流，而不必手动处理掩码或向量化的复杂性。然而，这种灵活性是有代价的：当warp内的线程出现分支分歧时，硬件必须串行执行各个分支路径，导致有效吞吐量下降。

硬件实现机制：SIMT的硬件实现依赖于以下几个关键组件：

独立程序计数器：每个线程维护自己的PC值，记录下一条要执行的指令地址。这使得线程可以独立地跟随不同的控制流路径。

活动掩码（Activity Mask）：warp调度器维护一个活动掩码，标识warp内哪些线程当前是活跃的（即需要执行当前指令）。当所有线程走相同路径时，活动掩码全为1；当出现分支分歧时，活动掩码只标记需要执行当前分支的线程。

栈式分支管理：GPU硬件使用栈结构来管理嵌套分支。当遇到分支指令时，硬件将当前的活动掩码压栈，然后根据分支条件更新活动掩码。分支结束后，从栈中恢复之前的活动掩码。

收敛机制：为了确保分支后的正确执行，硬件需要确保所有线程在分支汇合点重新收敛。这通常通过在分支末尾插入同步屏障来实现。

执行效率分析：SIMT的执行效率可以用以下公式近似表示：

$$Efficiency = \frac{Active\_Threads}{Total\_Threads\_in\_Warp}$$

在最理想的情况下（无分支分歧），效率为100%。如果warp内有一半线程走一个分支，另一半走另一个分支，则效率降为50%，因为两个分支需要串行执行。在最坏情况下（每个线程走不同的分支），效率仅为1/32 ≈ 3.125%。

需要注意的是，分支分歧并不是唯一影响SIMT效率的因素。内存访问模式、共享内存bank conflict、原子操作竞争等也会影响实际执行效率。但分支分歧是最直接、最容易理解的影响因素。

2.2 Warp：GPU并行执行的基本单位

Warp是NVIDIA GPU中线程组织的基本执行单位，理解warp的特性对于编写高效的CUDA代码至关重要。

Warp的定义与组成：一个warp由32个线程组成，这些线程在同一个block内连续分配。具体来说，对于一个block，线程索引0-31构成第一个warp，32-63构成第二个warp，依此类推。warp内的线程ID可以通过threadIdx.x % 32来计算。

选择32作为warp大小是经过深思熟虑的设计决策。较大的warp可以提高SIMD效率（更多的线程共享同一个指令解码器），但会增加分支分歧的概率和寄存器压力。较小的warp可以减少分支分歧的影响，但会增加硬件开销（每个线程需要独立的PC和状态）。32是一个经验上的平衡点，在大多数工作负载下表现良好。

Warp的执行特性：

锁步执行（Lock-step Execution）：warp内的所有活跃线程在同一时钟周期执行相同的指令。这意味着如果一条指令需要4个时钟周期完成，整个warp都需要等待4个周期。

快速上下文切换：由于线程状态主要保存在寄存器中，warp之间的切换几乎没有开销。当一个warp等待内存访问时，调度器可以立即切换到另一个就绪的warp，从而隐藏内存延迟。

隐式同步：warp内的天然是同步的，不需要显式的同步原语。程序员可以利用这一特性实现warp级别的原语，如warp shuffle、warp vote等。

分支序列化：如前所述，当warp内出现分支分歧时，硬件会串行执行各个分支路径。在执行某个分支时，不走该分支的线程被禁用（masked off），不产生任何计算结果。

Warp调度策略：现代GPU的SM包含多个warp调度器，每个调度器可以在每个时钟周期选择一个就绪的warp发射指令。调度器的选择策略通常基于优先级和就绪状态：

• 就绪优先：优先选择所有操作数都已就绪的warp
• 公平调度：避免某个warp长期得不到执行
• 年龄优先：优先选择等待时间最长的warp

不同的GPU架构可能采用不同的调度策略，但核心目标都是最大化执行单元的利用率。

Warp级别的优化原语：CUDA提供了一系列warp级别的内置函数，允许线程在warp内高效地交换数据和进行归约操作：

• __shfl_sync()：warp shuffle，允许线程从warp内的其他线程读取寄存器值
• __ballot_sync()：warp vote，收集warp内所有线程的条件判断结果
• __activemask()：返回当前活跃的线程掩码

这些原语利用了warp的硬件特性，比使用共享内存进行线程间通信更高效。例如，warp shuffle可以直接通过寄存器文件交换数据，避免了共享内存的加载-存储开销。

2.3 Streaming Multiprocessor (SM) 微架构

SM（Streaming Multiprocessor）是GPU的核心计算单元，包含了执行CUDA kernel所需的所有资源。理解SM的微架构对于性能优化至关重要。

SM的核心组件：

Warp调度器（Warp Scheduler）：负责选择就绪的warp并发射指令。现代SM通常包含多个独立的调度器，每个调度器可以独立地选择warp和发射指令。例如，Ampere架构的SM有4个调度器，每个时钟周期最多可以发射4条指令。

指令分发单元（Dispatch Unit）：将调度器选择的指令分发给相应的执行单元。指令分发需要考虑执行单元的可用性和操作数的就绪状态。

执行单元（Execution Units）：

- CUDA Core：执行整数和单精度浮点运算

- Tensor Core：执行矩阵乘法加速运算（Volta及以后架构）

- SFU（Special Function Unit）：执行 transcendental 函数（sin, cos, exp, log等）

- LD/ST Unit：执行内存加载和存储操作

- DP Unit：执行双精度浮点运算（部分架构中独立于CUDA Core）

寄存器文件（Register File）：SM内所有线程共享的大型寄存器堆。寄存器文件被划分为多个bank，以支持高带宽的并行访问。例如，Ampere SM的寄存器文件包含65536个32位寄存器。

共享内存/L1缓存：可配置的片上存储器，既可以用作程序员显式管理的共享内存，也可以用作自动缓存的L1缓存。在Ampere架构中，共享内存和L1缓存共享192KB的物理存储空间，可以灵活配置比例。

常量缓存（Constant Cache）：只读缓存，用于加速常量内存的访问。常量缓存在warp广播机制下特别高效，因为warp内所有线程读取相同的常量地址时，只需要一次内存事务。

纹理单元（Texture Unit）：专门用于纹理采样的硬件单元，支持过滤、地址模式等图形操作。在通用计算中，纹理单元也可以用于特定的数据访问模式。

SM的资源限制：

每个SM的资源是有限的，这限制了可以同时驻留的block数量和线程数量。主要的资源限制包括：

• 最大线程数：每个SM最多支持的活跃线程数（如Ampere为1536或2048）
• 最大Block数：每个SM最多支持的活跃block数（如Ampere为16或32）
• 寄存器数量：每个SM的寄存器文件总大小，限制了每个线程可以使用的寄存器数量
• 共享内存容量：每个SM的共享内存总量，限制了每个block可以分配的共享内存大小

这些限制共同决定了SM的occupancy（占用率），即活跃warp数量与最大支持warp数量的比值。Occupancy是影响性能的重要因素，但并非越高越好，需要与其他优化目标权衡。

SM的多世代演进：

每一代架构都在保持SIMT核心模型的同时，针对特定应用领域进行了优化。Volta引入的独立线程调度（Independent Thread Scheduling）尤其值得注意，它允许更细粒度的线程同步，使得warp级别的协作更加灵活。

2.4 分支处理与发散管理

分支分歧（Branch Divergence）是SIMT执行模型中最常见的性能陷阱之一。深入理解分支处理的硬件机制，有助于编写更高效的代码。

分支分歧的产生：当warp内的线程需要根据不同的条件执行不同的代码路径时，就产生了分支分歧。例如：

if (threadIdx.x % 2 == 0) {
    // 偶数线程执行
    a = b + c;
} else {
    // 奇数线程执行
    a = b - c;
}

在这个例子中，warp内的一半线程（偶数索引）执行加法，另一半（奇数索引）执行减法。由于warp是锁步执行的，硬件必须首先执行加法分支（禁用奇数线程），然后执行减法分支（禁用偶数线程）。这使得有效吞吐量减半。

硬件处理机制：GPU使用栈式分支管理来处理嵌套分支：

分支指令：当遇到条件分支时，硬件评估所有线程的条件，生成活动掩码。

压栈：将当前的活动掩码和返回地址压入分支栈。

更新掩码：根据分支条件更新活动掩码，只保留需要执行当前分支的线程。

执行分支：使用更新后的活动掩码执行分支体内的指令。禁用的线程不执行任何操作，不消耗执行单元资源，但仍然占用指令发射槽。

弹栈与收敛：分支结束后，从栈中恢复之前的活动掩码，确保所有线程在汇合点重新收敛。

对于嵌套分支，这个过程会递归进行。分支栈的深度限制了可以处理的嵌套分支层数，通常为24-32层。

分支优化的策略：

数据重排：通过重新组织数据，使得相邻线程倾向于走相同的分支路径。例如，将需要执行相同操作的元素聚集在一起。

分支消除：使用算术运算代替条件分支。例如，将if (cond) a = b; else a = c;改写为a = cond ? b : c;或使用掩码：a = mask * b + (1 - mask) * c;。

函数指针/跳转表：对于多路分支，使用查找表代替嵌套if-else，减少分支层级。

循环展开：对于循环内的分支，展开循环可以将分支条件提升为编译时常量，允许编译器优化掉不必要的分支。

Warp级投票：使用__any_sync()或__all_sync()检查warp内是否有线程满足条件，如果所有线程走相同路径，则可以跳过分支处理。

独立线程调度（ITS）：Volta架构引入的独立线程调度是对传统SIMT模型的重要改进。在ITS之前，warp内的线程严格锁步执行，同步只能在warp级别进行。ITS允许线程在更细粒度上独立推进，使得线程间的同步更加灵活。

ITS的主要优势：

• 支持线程级别的同步原语（如__syncwarp()）
• 减少了不必要的序列化
• 使得生产者-消费者模式在warp内更加高效

然而，ITS也增加了编程的复杂性，程序员需要更加小心地处理线程间的依赖关系，避免竞态条件。

2.5 指令流水线与执行单元

GPU的指令流水线设计与CPU有显著差异，反映了两者不同的设计目标。

流水线深度：GPU的流水线相对较浅，通常为10-20级，而现代CPU的流水线可达30-40级。较浅的流水线降低了分支预测失败的惩罚，但限制了最高时钟频率。这对于GPU来说是合理的选择，因为GPU依靠大量线程的并行执行来获得高吞吐量，而不是依靠高时钟频率。

顺序执行：GPU采用顺序执行模型，指令按照程序顺序依次发射和执行。这与CPU的乱序执行形成鲜明对比。顺序执行简化了硬件设计，减少了控制逻辑的面积和功耗开销。内存延迟通过快速线程切换来隐藏，而不是通过复杂的乱序执行机制。

延迟隐藏机制：GPU隐藏内存延迟的主要机制是上下文切换。当一个warp等待内存访问时，调度器立即切换到另一个就绪的warp。由于线程状态保存在寄存器中，切换开销极小。为了有效地隐藏延迟，需要有足够的活跃warp。这就是为什么occupancy对性能如此重要。

执行单元的组织：

GPU的执行单元分为以下几类：

INT32单元：执行32位整数运算。在Ampere架构中，INT32单元与FP32单元分离，可以并行执行。

FP32单元（CUDA Core）：执行单精度浮点运算。这是GPU最主要的计算单元，数量最多。

FP64单元：执行双精度浮点运算。在某些架构中（如GP100），FP64单元有独立的硬件；在其他架构中，FP64通过FP32单元模拟，性能较低。

Tensor Core：专门用于矩阵乘法加速的单元。Volta引入的Tensor Core可以执行4x4矩阵的乘加运算，支持FP16、INT8、INT4等低精度格式。Ampere的Tensor Core进一步支持TF32格式，在保持FP32精度的同时获得接近FP16的性能。

SFU（Special Function Unit）：执行超越函数（sin、cos、exp、log、sqrt、rsqrt等）。SFU的计算通常需要多个时钟周期，精度略低于标准浮点单元。

LD/ST单元：执行内存加载和存储操作。LD/ST单元负责将数据从内存子系统传送到寄存器，或反之。

指令吞吐与延迟：

不同类型的指令有不同的吞吐量和延迟特性：

吞吐量表示每个时钟周期可以启动的指令数量，延迟表示指令完成所需的时钟周期数。高吞吐量和低延迟是理想情况，但在实际硬件中往往需要权衡。

指令级并行（ILP）：虽然GPU主要依赖线程级并行（TLP），但指令级并行也很重要。编译器会尝试重新排序指令，使得长延迟指令（如内存加载）的结果在使用前有足够的时间内完成。程序员也可以通过软件预取等技术帮助编译器提高ILP。

2.6 寄存器文件与共享内存子系统

寄存器文件和共享内存是GPU片上存储层次的核心组件，它们的性能直接影响kernel的执行效率。

寄存器文件：

GPU的寄存器文件是一个大型的、多端口的寄存器堆，为SM内所有活跃线程提供寄存器存储。与CPU不同，GPU的寄存器分配是在编译时静态确定的，而不是在运行时动态分配。

关键特性：

大容量：Ampere SM的寄存器文件包含65536个32位寄存器，足以支持数千个线程同时驻留。

零开销切换：由于每个线程的寄存器状态在物理上是隔离的（通过寄存器索引偏移），warp切换不需要保存和恢复寄存器状态，切换开销为零。

静态分配：编译器在编译时确定每个kernel需要的寄存器数量。过多的寄存器使用会限制occupancy，因为SM的寄存器总数是有限的。

溢出处理：如果kernel需要的寄存器超过可用数量，编译器会将部分变量溢出到本地内存（实际上是全局内存的一部分），这会显著降低性能。应该尽量避免寄存器溢出。

寄存器压力与Occupancy权衡：

每个线程使用的寄存器数量直接影响SM可以容纳的线程数量。假设SM有65536个寄存器，如果每个线程使用32个寄存器，则SM最多可以容纳65536/32 = 2048个线程。如果每个线程使用64个寄存器，则最多只能容纳1024个线程。

CUDA提供了__launch_bounds__编译指示，允许程序员指定预期的block大小和最大寄存器使用量，帮助编译器做出更好的寄存器分配决策。也可以使用maxrregcount编译器选项限制寄存器的最大使用量，强制编译器将部分变量溢出，以换取更高的occupancy。这需要仔细的性能测试来确定最佳平衡点。

共享内存：

共享内存是GPU特有的高速片上存储器，具有以下特点：

低延迟：共享内存的访问延迟约为1-2个时钟周期，与寄存器相当，远快于全局内存。

高带宽：共享内存的带宽远高于全局内存，可以支持多个线程同时访问。

可编程管理：与缓存不同，共享内存由程序员显式管理。程序员负责数据的加载、存储和同步。

线程间通信：同一block内的线程可以通过共享内存交换数据，实现协作。

Bank结构：共享内存被划分为多个bank（通常32个），每个bank每个时钟周期可以服务一个32位字的访问。如果多个线程同时访问同一个bank的不同地址，就会发生bank conflict，访问需要串行化。

Bank Conflict详解：

共享内存的32个bank分别对应warp的32个线程。理想情况下，warp内的32个线程访问32个不同bank的地址，可以在一个时钟周期内完成所有访问。但如果两个或多个线程访问同一个bank的不同地址，就会发生bank conflict。

例如，假设共享内存数组float shared[32]，如果thread 0访问shared[0]，thread 1访问shared[1]，...，thread 31访问shared[31]，则没有bank conflict，因为每个线程访问不同的bank。

但如果thread 0访问shared[0]，thread 1访问shared[2]，thread 2访问shared[4]，...，则thread 0和thread 16都访问bank 0（因为bank索引 = 地址索引 % 32），发生2-way bank conflict，访问需要2个时钟周期。

避免bank conflict的策略：

填充（Padding）：在数组第二维加1，改变地址映射，避免多个线程访问同一bank。例如，float shared[32][33]而不是float shared[32][32]。

数据重排：重新组织数据结构，使得相邻线程访问相邻的地址。

广播机制：如果warp内所有线程访问相同的地址，共享内存支持广播，不会产生bank conflict。

共享内存与L1缓存的统一：

在较新的GPU架构中，共享内存和L1缓存共享同一块物理存储空间。程序员可以通过CUDA API配置两者的比例。例如，在Ampere中，可以配置为112KB共享内存 + 80KB L1缓存，或80KB共享内存 + 112KB L1缓存等。这种灵活性允许根据应用特点优化存储层次的使用。

三、程序员的视角

3.1 在实际项目中的应用

理解GPU架构在实际项目中有着广泛的应用：

性能调优：通过分析kernel的occupancy、寄存器使用、共享内存使用等指标，可以识别性能瓶颈并针对性优化。例如，如果occupancy过低是因为寄存器使用过多，可以考虑减少每线程的寄存器使用或调整block大小。

算法选择：不同的算法在GPU上的表现差异很大。例如，在排序算法中，bitonic sort的比较模式规则，适合GPU并行；而快速排序的递归和不规则访问模式在GPU上效率较低。理解架构特点有助于选择适合的算法变体。

库开发：高性能GPU库（如cuBLAS、cuFFT）的开发需要对架构有深入的理解。库开发者需要手工优化汇编代码，精细控制寄存器分配和指令调度，以达到接近理论峰值的性能。

新兴应用适配：对于深度学习、图计算等新兴应用，理解架构有助于设计专用的数据结构和访问模式。例如，在图神经网络中，节点特征的聚合操作可以通过共享内存优化，减少全局内存访问。

3.2 常见陷阱与误区

过度追求高Occupancy：Occupancy高并不一定意味着性能好。有时降低occupancy可以增加每线程可用的寄存器和共享内存，从而改善数据局部性或允许更复杂的计算。应该以实际性能测量为准，而不是盲目追求高occupancy。

忽视Bank Conflict：共享内存的bank conflict可能导致严重的性能下降，但初学者常常忽视这一点。使用Nsight Compute工具可以检测bank conflict，指导优化。

错误理解Warp同步：在Volta之前的架构，warp内天然同步，不需要显式同步。但Volta引入ITS后，线程可以独立推进，需要使用__syncwarp()确保warp内同步。忽视这一点可能导致竞态条件。

寄存器溢出：当kernel使用过多寄存器时，编译器会将部分变量溢出到本地内存（全局内存），导致性能急剧下降。应该监控寄存器使用量，必要时限制最大寄存器数以保持合理的occupancy。

假设固定的SM数量：不同GPU型号的SM数量差异很大。编写可移植的代码时，不应该硬编码SM数量，而应该使用cudaGetDeviceProperties()查询设备属性。

3.3 与其他层级的联系

与编译器的联系：CUDA编译器（nvcc）负责将高级语言代码转换为GPU机器码。编译器的优化决策（如寄存器分配、指令调度、循环展开）直接影响最终性能。理解编译器的工作原理有助于编写更容易优化的代码。

与运行时的联系：CUDA运行时负责kernel的启动、内存管理、同步等操作。Kernel启动是异步的，运行时将kernel放入命令队列，由驱动调度执行。理解这一机制有助于设计重叠计算和数据传输的策略。

与硬件的联系：最终的执行效果取决于硬件的具体实现。不同架构的GPU在指令延迟、带宽、缓存大小等方面有差异。性能敏感的代码应该针对目标硬件进行调优。

四、代码实战

4.1 Warp Shuffle示例：高效的Warp级归约

Warp shuffle允许线程直接从其他线程的寄存器读取数据，无需通过共享内存，是实现warp级归约的高效方法。

#include <cuda_runtime.h>
#include <stdio.h>

// Warp-level reduction using shuffle instructions
__device__ float warpReduceSum(float val) {
    // Lane ID within warp
    int laneId = threadIdx.x % 32;
    
    // Butterfly reduction pattern
    #pragma unroll
    for (int offset = 16; offset > 0; offset /= 2) {
        val += __shfl_down_sync(0xFFFFFFFF, val, offset);
    }
    
    return val;
}

// Block-level reduction using warp shuffle + shared memory
__global__ void blockReduceSum(const float* input, float* output, int N) {
    extern __shared__ float shared[];
    
    int tid = threadIdx.x;
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    // Load value from global memory
    float val = (idx < N) ? input[idx] : 0.0f;
    
    // Warp-level reduction
    val = warpReduceSum(val);
    
    // Store warp result to shared memory
    // Only first thread in each warp writes
    if (laneId == 0) {
        shared[tid / 32] = val;
    }
    
    __syncthreads();
    
    // Final reduction in first warp
    int numWarps = (blockDim.x + 31) / 32;
    if (tid < numWarps) {
        val = shared[tid];
        val = warpReduceSum(val);
        
        // First thread writes final result
        if (tid == 0) {
            atomicAdd(output, val);
        }
    }
}

// Host function to test block reduction
void testBlockReduction(int N) {
    size_t size = N * sizeof(float);
    
    // Allocate and initialize host data
    float *h_input = (float*)malloc(size);
    float h_output = 0.0f;
    
    for (int i = 0; i < N; i++) {
        h_input[i] = 1.0f;
    }
    
    // Allocate device memory
    float *d_input, *d_output;
    cudaMalloc((void**)&d_input, size);
    cudaMalloc((void**)&d_output, sizeof(float));
    
    cudaMemcpy(d_input, h_input, size, cudaMemcpyHostToDevice);
    cudaMemset(d_output, 0, sizeof(float));
    
    // Launch kernel
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    int sharedMemSize = (threadsPerBlock / 32) * sizeof(float);
    
    blockReduceSum<<<blocksPerGrid, threadsPerBlock, sharedMemSize>>>(
        d_input, d_output, N);
    
    // Copy result back
    cudaMemcpy(&h_output, d_output, sizeof(float), cudaMemcpyDeviceToHost);
    
    printf("Reduction result: %f (expected: %d)\n", h_output, N);
    
    // Cleanup
    cudaFree(d_input);
    cudaFree(d_output);
    free(h_input);
}

int main() {
    testBlockReduction(1000000);
    return 0;
}

这个示例展示了：

使用__shfl_down_sync实现warp级归约，无需共享内存

结合warp shuffle和共享内存实现block级归约

使用atomicAdd汇总各block的结果

4.2 分支分歧性能测试

通过实验量化分支分歧对性能的影响。

#include <cuda_runtime.h>
#include <stdio.h>

// Kernel with no branch divergence
__global__ void noBranchKernel(float* data, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        data[idx] = sqrtf(fabs(data[idx]));
    }
}

// Kernel with branch divergence (alternating threads)
__global__ void divergentKernel(float* data, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        if (idx % 2 == 0) {
            data[idx] = sqrtf(fabs(data[idx]));
        } else {
            data[idx] = data[idx] * data[idx];
        }
    }
}

// Kernel with branch divergence (by warp)
__global__ void warpAlignedBranchKernel(float* data, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        // All threads in a warp take the same path
        if ((idx / 32) % 2 == 0) {
            data[idx] = sqrtf(fabs(data[idx]));
        } else {
            data[idx] = data[idx] * data[idx];
        }
    }
}

// Benchmark utility
void benchmarkKernel(void (*kernel)(float*, int), const char* name, 
                     float* d_data, int N, int iterations) {
    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    
    // Warm-up
    for (int i = 0; i < 5; i++) {
        kernel<<<(N + 255) / 256, 256>>>(d_data, N);
    }
    cudaDeviceSynchronize();
    
    // Measure
    cudaEventRecord(start);
    for (int i = 0; i < iterations; i++) {
        kernel<<<(N + 255) / 256, 256>>>(d_data, N);
    }
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);
    
    float ms = 0;
    cudaEventElapsedTime(&ms, start, stop);
    
    printf("%-30s: %.3f ms (%.3f ms/iter)\n", name, ms, ms / iterations);
    
    cudaEventDestroy(start);
    cudaEventDestroy(stop);
}

int main() {
    int N = 10000000;
    size_t size = N * sizeof(float);
    
    float *h_data = (float*)malloc(size);
    for (int i = 0; i < N; i++) {
        h_data[i] = (float)i - N / 2;
    }
    
    float *d_data;
    cudaMalloc((void**)&d_data, size);
    cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice);
    
    int iterations = 100;
    
    printf("=== Branch Divergence Benchmark ===\n");
    printf("Data size: %d elements\n", N);
    printf("Iterations: %d\n\n", iterations);
    
    benchmarkKernel(noBranchKernel, "No Branch", d_data, N, iterations);
    benchmarkKernel(divergentKernel, "Divergent (alt threads)", d_data, N, iterations);
    benchmarkKernel(warpAlignedBranchKernel, "Divergent (by warp)", d_data, N, iterations);
    
    cudaFree(d_data);
    free(h_data);
    
    return 0;
}

4.3 共享内存Bank Conflict演示

演示bank conflict的产生和避免方法。

#include <cuda_runtime.h>
#include <stdio.h>

#define TILE_SIZE 32

// Kernel with bank conflicts
__global__ void bankConflictKernel(float* output, int N) {
    __shared__ float shared[TILE_SIZE][TILE_SIZE];
    
    int tx = threadIdx.x;
    int ty = threadIdx.y;
    
    // Each thread accesses column = tx, row = (tx + ty) % TILE_SIZE
    // This causes bank conflicts because multiple threads access same bank
    int row = (tx + ty) % TILE_SIZE;
    shared[row][tx] = (float)(row * TILE_SIZE + tx);
    
    __syncthreads();
    
    if (tx == 0 && ty == 0) {
        output[0] = shared[0][0];
    }
}

// Kernel without bank conflicts (using padding)
__global__ void noBankConflictKernel(float* output, int N) {
    // Add padding to avoid bank conflicts
    __shared__ float shared[TILE_SIZE][TILE_SIZE + 1];
    
    int tx = threadIdx.x;
    int ty = threadIdx.y;
    
    int row = (tx + ty) % TILE_SIZE;
    shared[row][tx] = (float)(row * TILE_SIZE + tx);
    
    __syncthreads();
    
    if (tx == 0 && ty == 0) {
        output[0] = shared[0][0];
    }
}

// Benchmark
void benchmarkBankConflict() {
    float *d_output;
    cudaMalloc((void**)&d_output, sizeof(float));
    
    dim3 blockSize(TILE_SIZE, TILE_SIZE);
    dim3 gridSize(1, 1);
    
    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    
    int iterations = 1000;
    
    // Warm-up
    for (int i = 0; i < 10; i++) {
        bankConflictKernel<<<gridSize, blockSize>>>(d_output, 0);
    }
    cudaDeviceSynchronize();
    
    // Measure bank conflict version
    cudaEventRecord(start);
    for (int i = 0; i < iterations; i++) {
        bankConflictKernel<<<gridSize, blockSize>>>(d_output, 0);
    }
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);
    
    float ms1 = 0;
    cudaEventElapsedTime(&ms1, start, stop);
    
    // Measure no bank conflict version
    cudaEventRecord(start);
    for (int i = 0; i < iterations; i++) {
        noBankConflictKernel<<<gridSize, blockSize>>>(d_output, 0);
    }
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);
    
    float ms2 = 0;
    cudaEventElapsedTime(&ms2, start, stop);
    
    printf("=== Bank Conflict Benchmark ===\n");
    printf("With bank conflicts:    %.3f ms (%.3f ms/iter)\n", ms1, ms1 / iterations);
    printf("Without bank conflicts: %.3f ms (%.3f ms/iter)\n", ms2, ms2 / iterations);
    printf("Speedup: %.2fx\n", ms1 / ms2);
    
    cudaFree(d_output);
    cudaEventDestroy(start);
    cudaEventDestroy(stop);
}

int main() {
    benchmarkBankConflict();
    return 0;
}

五、深入思考

5.1 历史背景

GPU架构的演进反映了计算需求的变迁。早期的GPU是固定的功能图形管线，程序员无法编程。2001年，NVIDIA GeForce 3引入了可编程顶点和像素着色器，但仍受限于图形操作。2006年，Tesla架构实现了统一着色器架构，所有执行单元可以执行任意类型的着色器代码，这是GPGPU的基础。

Fermi架构（2010）引入了真正的缓存层次和ECC内存支持，使GPU更适合科学计算。Kepler（2012）引入了动态并行，允许kernel启动子kernel。Maxwell（2014）专注于能效优化，为移动和嵌入式市场铺平道路。

Volta（2017）是面向AI的关键架构，引入了Tensor Core和独立线程调度。Tensor Core专门加速矩阵乘法，成为深度学习训练的基石。独立线程调度打破了严格的warp锁步执行，提供了更灵活的同步原语。

Ampere（2020）进一步增强了Tensor Core，引入了TF32格式和结构化稀疏性支持。Hopper（2022）针对Transformer模型优化，引入了Transformer Engine和DPX指令。

5.2 前沿方向

Chiplet集成：未来的GPU可能采用Chiplet设计，将计算、内存、I/O等功能分解为独立的小芯片，通过先进封装技术集成。这可以提高良率、降低成本，并允许混合工艺节点。

光互联：随着带宽需求的不断增长，电信号传输面临瓶颈。硅光技术利用光子进行芯片间通信，可以提供更高的带宽和更低的功耗。

近存计算：将计算单元靠近或嵌入存储器内部，减少数据搬运开销。HBM-PIM是这一方向的早期尝试。

专用加速器：通用GPU与专用加速器（如TPU、NPU）的界限正在模糊。未来的GPU可能集成更多专用单元，针对特定工作负载（如图处理、稀疏计算）进行优化。

5.3 延伸阅读

NVIDIA Volta Architecture Whitepaper

NVIDIA Ampere Architecture Whitepaper

NVIDIA Hopper Architecture Whitepaper

"Inside the NVIDIA Volta GPU Architecture", IEEE Micro, 2018

CUDA C Programming Guide - Occupancy Calculator

六、实验笔记

实验1：Warp Shuffle vs 共享内存归约性能对比

目的：比较warp shuffle和共享内存在归约操作中的性能差异。

步骤：

实现两种版本的block级归约：一种使用warp shuffle，一种仅使用共享内存

测试不同block大小下的执行时间

分析两种方法的寄存器使用和共享内存使用差异

预期：warp shuffle版本应该更快，因为它减少了共享内存访问次数，且shuffle指令通过寄存器文件直接交换数据，延迟更低。

实验2：Occupancy对性能的影响

目的：研究occupancy与kernel性能的关系。

步骤：

编写一个简单的计算密集型kernel

使用__launch_bounds__限制最大寄存器数，人为降低occupancy

测量不同occupancy下的执行时间

重复上述步骤，但使用内存密集型kernel

预期：对于计算密集型kernel，高occupancy有助于隐藏指令延迟；对于内存密集型kernel，低occupancy可能允许更多寄存器用于数据缓存，反而提高性能。

七、本章小结

本章深入探讨了GPU架构的核心要素，重点分析了SIMT执行模型的工作原理。我们从warp的概念出发，理解了GPU如何将线程组织为基本的执行单位，以及warp调度器如何实现零开销的上下文切换。通过对SM微架构的详细剖析，我们认识了warp调度器、执行单元、寄存器文件、共享内存等关键组件，以及它们如何协同工作以实现高吞吐量的并行计算。

分支分歧是SIMT模型中最常见的性能陷阱，我们通过硬件机制的分析和优化策略的讨论，学习了如何编写分支友好的代码。Warp shuffle等warp级原语为我们提供了高效的线程间通信手段，超越了传统的共享内存方式。

寄存器文件和共享内存子系统是GPU片上存储层次的核心，它们的容量和带宽直接制约着kernel的性能。我们讨论了寄存器压力与occupancy的权衡，以及共享内存bank conflict的产生机制和避免方法。

通过三个代码示例，我们实践了warp级归约、分支分歧测试和bank conflict优化，将理论知识转化为实际的编程技能。最后，我们回顾了GPU架构的历史演进，展望了Chiplet、光互联、近存计算等前沿方向。

理解GPU架构是成为优秀CUDA程序员的必经之路。只有深入理解硬件的工作机制，才能做出明智的设计决策，编写出真正高效的并行代码。

八、习题

Warp计算：一个CUDA block包含256个线程。计算这个block包含多少个warp？如果SM最多支持1536个线程，且没有其他限制，这个SM最多可以同时驻留多少个这样的block？

分支分析：给定kernel代码if (threadIdx.x < 100) { ... } else { ... }，block大小为256。分析有多少个warp会出现分支分歧？对于出现分歧的warp，有效执行效率是多少？

寄存器限制：假设SM有65536个32位寄存器，每个线程使用40个寄存器。计算SM最多可以容纳多少个线程？如果block大小为256，SM最多可以驻留多少个block？此时的occupancy是多少（假设SM最大支持1536线程）？

Bank Conflict分析：共享内存数组定义为float shared[32][32]。如果thread (tx, ty)访问shared[ty][tx]，是否存在bank conflict？如果改为访问shared[tx][ty]呢？如果将数组改为float shared[32][33]，再次访问shared[ty][tx]，情况如何？

Shuffle实现：使用__shfl_xor_sync实现一个warp级的最大值归约函数。解释蝴蝶归约模式的工作原理。

架构比较：比较Volta和Ampere架构在Tensor Core方面的主要差异。TF32格式的特点是什么？它在精度和性能之间如何权衡？

延迟隐藏：解释GPU如何通过线程切换隐藏内存延迟。如果一个内存加载需要400个时钟周期完成，理论上需要多少个活跃warp才能完全隐藏这个延迟？假设每个warp每条指令发射后需要等待结果才能发射下一条。

优化策略：对于一个寄存器使用量很高的kernel，occupancy只有25%。列出至少三种可能的优化策略来提高性能，并说明每种策略的适用场景和潜在风险。

ITS影响：Volta引入的独立线程调度（ITS）对编程模型有什么影响？在ITS之前，warp内的线程天然同步；ITS之后，为什么需要使用__syncwarp()？给出一个因忽视__syncwarp()而导致错误的代码示例。

综合设计：设计一个高效的并行直方图计算kernel。说明如何利用共享内存减少全局内存原子操作，如何处理bin冲突，以及如何利用warp级原语进一步优化。