揭秘NVIDIA GT200微架构：通过微基准测试发现未公开的存储层级与同步机制

本文不仅验证了CUDA编程指南[1]中记录的部分硬件特性，还揭示了一系列未在文档中公开的硬件结构，例如_控制流机制、缓存与TLB层级_。此外，在某些场景下，我们的发现与文档描述的特性存在差异（例如纹理缓存和常量缓存的行为）。

本文的核心价值在于介绍了一套用于GPU架构分析的方法论。我们相信，这些方法对于分析其他类型的GPU架构以及验证类GPU性能模型都将有所助益。

本研究的最终目标是更深入地理解GPU硬件，从而充分挖掘其性能潜力。

英伟达GT200 GPU芯片的晶圆核心微观结构（显微拍摄）。GT200采用Tesla 2.0架构，由台积电以65纳米工艺制造。其芯片尺寸为576平方毫米，集成了14亿个晶体管，是一款规模庞大的芯片。GT200支持DirectX 11.1（功能级别10_0）。在GPU计算方面，它支持OpenCL 1.1和CUDA 1.3。该芯片拥有240个流处理器（着色单元）、80个纹理映射单元和32个光栅操作单元。

• Demystifying GPU Microarchitecture through Microbenchmarking
• https://www.stuffedcow.net/files/gpuarch-ispass2010.pdf
• Decuda：https://github.com/laanwj/decuda

《Demystifying GPU Microarchitecture through Microbenchmarking》这篇文章聚焦于NVIDIA GT200（GTX280）GPU，旨在通过微基准测试来揭示其微架构的底层细节。

尽管GPU在非图形计算中能实现数量级的性能提升，但CUDA的类C语言抽象掩盖了底层硬件特征，同时厂商提供的文档也较为有限。为此，研究团队开发了一套专用的微基准测试套件，重点探究影响性能的两大关键模块：算术处理核心与存储层次结构。

研究验证了CUDA编程指南中的部分性能特征，并深入剖析了分支发散与屏障同步的工作机制，发现了可能导致死锁的非直观分支代码序列。通过使用Decuda工具分析原生指令，团队精确测量了SP（流处理器）、SFU（特殊函数单元）、DPU（双精度单元）三类算术流水线的延迟与吞吐量，明确了各类运算的执行单元映射规律。

在存储层次方面，研究揭示了多级缓存与TLB结构：指令内存、常量内存、纹理内存均存在多级缓存，全局内存与纹理内存配备了两级TLB，并且各类存储的缓存大小、关联度、行大小等关键参数均被精准量化。此外，还发现了syncthreads()以warp为粒度进行同步、寄存器文件存在64个逻辑bank等未公开特征。

该研究不仅填补了GT200架构细节的空白，为程序优化、死锁规避及周期精确建模提供了数据支撑，其测试方法也为其他GPU及类GPU系统的架构分析提供了参考，有助于充分挖掘GPU硬件潜力。

本文目录

• 关键问题
  • 问题 1：厂商刻意隐瞒 GT200 核心架构细节的核心诉求是什么？是为了维持硬件技术壁垒，还是故意通过信息不对称限制开发者对 GPU 性能的极致挖掘？
  • 问题 2：syncthreads()实际同步粒度与文档描述不符，是否说明厂商在编程模型抽象时存在设计缺陷？为何厂商拒绝修正误导性描述？
• 一、引言
• 二、背景知识：GPU 架构与编程模型
  • A. GPU 架构
  • B. CUDA 软件编程接口
• 三、测量方法
  • A. 微基准测试方法
  • B. 从延迟图推断缓存特征
• 四、测试与结果
  • A. 时钟开销与特性
  • B. 运算流水线
  • C. 控制流：warp内线程的发散与收敛
  • D. 屏障同步
  • E. 寄存器文件
  • F. 共享内存
  • G. 全局内存
  • H. 纹理内存
  • I. 内存地址转换
  • J. 常量内存
  • K. 指令供给
• 五、相关工作
• 六、总结与结论
• 参考文献

关键问题

问题 1：厂商刻意隐瞒 GT200 核心架构细节的核心诉求是什么？是为了维持硬件技术壁垒，还是故意通过信息不对称限制开发者对 GPU 性能的极致挖掘？

论文揭示的 GT200 常量内存三级缓存、全局/纹理内存两级 TLB 等核心架构细节均未出现在厂商官方文档中，而这些细节直接决定程序优化方向与死锁规避效果，厂商刻意隐瞒这些关键信息的核心诉求是什么？ 是为了维持硬件技术壁垒，还是 故意通过信息不对称限制开发者对 GPU 性能的极致挖掘？

从论文披露的信息及行业逻辑来看，厂商隐瞒核心架构细节的核心诉求是平衡技术壁垒、编程模型易用性与生态控制权，而非单纯限制开发者性能挖掘：

1. 维持技术壁垒是核心考量。论文揭示的常量内存三级缓存、全局/纹理内存两级 TLB 等结构，是 GT200 性能优化的关键硬件基础。这些细节属于厂商差异化竞争的核心技术资产，隐瞒可避免竞争对手快速模仿其缓存分层、TLB 调度等底层设计，保障产品市场优势。
2. 简化编程模型，降低开发者门槛。CUDA 的核心价值是提供类C语言的抽象接口，让开发者无需关注底层硬件细节即可快速开发。论文提到“CUDA 将 GPU 抽象为多核系统，隐藏了 TPC、SM 内部的硬件组织”，若公开过多缓存关联度、TLB 行大小等细节，会增加编程复杂度，违背其“降低非图形计算门槛”的产品定位。
3. 规避硬件实现依赖，保障生态兼容性。论文发现部分架构细节（如纹理 L1 缓存 5KB 大小、syncthreads()的 warp 粒度） 与文档描述存在偏差，若公开具体实现，开发者可能会针对这些细节编写“硬件绑定”的优化代码，导致后续 GPU 硬件迭代（如缓存大小、TLB 结构变更）时，大量现有代码失效，增加生态维护成本。

厂商并非刻意限制性能挖掘，而是通过“高层规则引导+底层细节隐藏”的方式，让开发者在其设定的框架内优化（如 CUDA 指南提示“线程块大小为 64 的倍数时性能更佳”，间接呼应论文发现的寄存器文件 64 逻辑 bank 设计），既保障开发效率，又维持对硬件迭代和生态走向的控制权。

问题 2：syncthreads()实际同步粒度与文档描述不符，是否说明厂商在编程模型抽象时存在设计缺陷？为何厂商拒绝修正误导性描述？

论文证实 syncthreads()实际以 warp 为同步粒度，与 CUDA 编程指南“线程级屏障”的描述严重不符，且此类文档与硬件行为的偏差还存在于纹理缓存延迟、分支收敛机制等关键模块。这是否说明厂商在编程模型抽象时存在设计缺陷？ 为何厂商拒绝通过官方渠道修正这些误导性描述，反而让开发者依赖逆向工程式的微基准测试来规避风险？

这一偏差并非编程模型的设计缺陷，而是硬件实现与高层抽象的合理权衡，厂商拒绝修正描述的核心原因是保障接口稳定性与生态兼容性：

1. 抽象与实现的权衡，而非设计缺陷。论文证实 syncthreads() 以 warp 为同步粒度，但文档描述为“线程级屏障”，本质是 CUDA 编程模型的“高层抽象简化”——将复杂的硬件同步逻辑（warp 调度、分支收敛）封装为统一接口，让开发者无需理解 SIMT 架构的底层细节即可编写代码。这种抽象符合 GPU“海量线程并行”的设计目标，若强制按线程级同步，会大幅增加硬件调度开销，降低并行效率，反而违背性能设计初衷。
2. 修正描述可能引发生态混乱。论文指出，依赖“线程级同步”假设的代码可能出现死锁，但大量现有 CUDA 程序已基于文档描述开发。 若修正描述，需开发者重新适配代码（如调整分支结构、同步时机），不仅增加开发成本，还可能导致部分旧代码失效。厂商选择维持现有描述，是优先保障生态兼容性而非绝对准确性。
3. 引导开发者依赖“规则而非实现”。论文提到厂商通过“模糊规则”（如“syncthreads()仅在条件对全线程块一致时使用”）间接规避风险，本质是引导开发者遵循编程模型的“契约式设计”，而非依赖具体硬件实现。这种方式可让厂商在后续硬件迭代中灵活调整同步机制（如改变 warp 大小、同步粒度），而无需修改上层接口，保障 CUDA 生态的长期稳定性。

综上，文档与硬件行为的偏差是“抽象简化”的必然结果，而非设计缺陷；厂商拒绝修正是基于生态兼容性和硬件迭代灵活性的理性选择，而非刻意误导。

一、引言

图形处理器（GPU）有望在某些非图形计算任务上，提供比传统处理器高出一个数量级以上的加速比。由于 GPU 通常以类 C 语言的抽象形式呈现（例如英伟达的 CUDA），除了制造商公开的文档外，人们对 GPU 架构的特征知之甚少。

本研究开发了一套微基准测试套件（microbenchmark suite，用于精确测量硬件特定特征的小型测试程序集合），并对英伟达 GT200（GTX280）GPU 的 CUDA 可见架构特征进行了测量。研究中测量了处理单元和存储层次结构的多种【未公开】特征。该分析揭示了影响程序性能和正确性的【未文档化特性】， 这些测量结果有助于改进该架构上的性能优化、分析和建模工作，并为理解该 GPU 开发过程中的决策提供了额外视角。

图形处理器（GPU）作为非图形计算处理器，其架构与传统顺序处理器（sequential processors，指按指令顺序依次执行的处理器，如常见的 CPU）不同。对于开发者以及 GPU 架构和编译器研究者而言，详细理解现代 GPU 设计的架构至关重要。

英伟达 G80 和 GT200 系列 GPU 能够通过类 C 语言的 CUDA 编程接口执行非图形计算。《CUDA 编程指南》以规则的形式提供了关于 GPU 性能特征的提示[1]。然而，这些规则有时较为模糊，且关于支撑这些规则的底层硬件组织信息十分有限。

本研究提出了一套针对架构特定部分的微基准测试套件。所呈现的测量结果聚焦于影响 GPU 性能的两个主要部分：

• 算术处理核心（arithmetic processing cores，执行算术运算的硬件单元）
• 以及为这些处理核心提供指令和数据的存储层次结构（memory hierarchies，指由不同速度、容量的存储设备构成的层级结构，如寄存器、缓存、内存等）。

要避免死锁、优化应用性能并实现周期精确的 GPU 性能建模，就需要精确理解处理核心和缓存层次结构。

具体而言，本研究包含以下工作：

• 验证《CUDA 编程指南》中列出的性能特征。
• 探究分支发散（branch divergence）和屏障同步（barrier synchronization）的详细功能。研究发现，某些非直觉的分支代码序列会导致死锁，而理解内部架构可避免此类情况。
• 测量存储缓存层次结构的结构和性能，包括转换后备缓冲器（TLB）层次结构、常量内存（constant memory）、纹理内存（texture memory）和指令内存缓存。
• 讨论所采用的测量技术，研究人员认为这些技术将有助于其他 GPU 及类 GPU 系统的分析与建模，并提升 GPU 性能建模与仿真的保真度。

本文其余部分结构如下：
* 第二节回顾 CUDA 计算模型；
* 第三节描述测量方法；
* 第四节呈现测量结果；
* 第五节回顾相关工作；
* 第六节总结研究发现。

二、背景知识：GPU 架构与编程模型

A. GPU 架构

CUDA 将 GPU 架构建模为多核系统。它将 GPU 的线程级并行性抽象为线程层次结构：线程网格（grid）包含线程块（block），线程块包含线程束（warp），线程束包含线程。

这些线程被映射到硬件资源层次结构上。线程块在线式多处理器（SM）内执行，如图 1 所示。

图 1：每个流式多处理器包含 8 个标量处理器

尽管编程模型中使用的是标量线程集合，但 SM 更接近一个 8 路向量处理器，该处理器对 32 路向量进行操作。

SM 内执行流的基本单位是线程束（warp）。
* 在 GT200 中，一个线程束由 32 个线程组成，并以 8 个为一组（sub-warp），在 8 个标量处理器上执行。

    32线程Warp = 4 × 8线程子组  
    每个子组 → 8个SP同时执行（1个批次）  
    4个子组 → 4个批次完成全部32线程的指令执行  

英伟达将这种架构安排称为单指令多线程（SIMT）——线程束中的每个线程同步执行相同指令，但允许每个线程独立分支。
* SM 包含算术单元，以及其他为线程块和线程私有的资源，例如每个线程块专属的共享内存（shared memory）和寄存器文件（register file）。

图 2：每个线程处理集群包含 3 个 SM

多个 SM 组成线程处理集群（TPC），其不仅包含多个 SM 以及共享资源的硬件单元，共享资源如缓存、纹理提取单元，其中大部分资源对程序员不可见，如图 2 所示。

图 3：包含 TPC 和内存 bank 的 GPU

从 CUDA 的视角来看，GPU 由线程处理集群（TPC）集合、互连网络（interconnection network）和存储系统（DRAM 内存控制器）构成，如图 3 所示。下面表 I 列出了英伟达公开的 GT200 参数。

表 I：英伟达公开的 GT200 参数

B. CUDA 软件编程接口

CUDA 通过类 C 语言并扩展抽象线程模型的方式呈现 GPU 架构。

在 CUDA 模型中，主机（host）CPU 代码可通过调用在 GPU 上执行的设备（device）函数来启动 GPU 核函数。由于 GPU 与主机 CPU 使用不同的指令集，CUDA 编译流程会使用不同的编译器分别编译 CPU 和 GPU 代码，以适配各自的指令集。
* GPU 代码首先被编译为 PTX“汇编代码”（PTX）。
* 随后被“汇编”为原生代码（native code）。
* 编译后的 CPU 代码和 GPU 代码会被合并为一个单一的“胖二进制文件”（fat binary）。

尽管 PTX 被描述为 GPU 代码的汇编级表示，但它仅是一种中间表示，无法用于详细分析或微基准测试。由于原生指令集与 PTX 不同，且编译器会对 PTX 代码进行优化，因此 PTX 代码无法准确反映实际执行的机器指令。

在大多数情况下，研究人员发现最有效的方式是使用 CUDA C 编写代码，然后通过 decuda 工具在原生代码层面验证生成的机器代码序列。

使用 decuda 主要是为了便捷，因为生成的指令序列可在原生 cubin 二进制文件（cubin）中得到验证。

三、测量方法

A. 微基准测试方法

为探究 GT200 架构，研究人员设计了微基准测试程序，以暴露待测量的每一项架构特征。研究结论通过分析微基准测试程序的执行时间得出。

在测量指令缓存参数时，使用 decuda 工具报告代码大小和位置，其结果与对编译后代码的分析一致。 研究人员还利用 decuda 查看 CUDA 编译器生成的原生指令序列，并分析用于处理分支发散和重汇聚（reconvergence）的代码。

微基准测试程序的一般结构为：GPU 核函数包含计时代码，计时代码围绕一段用于测试目标硬件的代码段（通常是展开的循环，执行多次）。
* 一个基准测试核函数会完整运行两次代码，第一次迭代的结果会被忽略，以避免冷指令缓存缺失的影响。
* 在所有情况下，核函数代码的大小都足够小，可放入 L1 指令缓存，以确保测试的是目标硬件特征的真实耗时，而非指令加载的耗时。
* 计时测量通过读取时钟寄存器（使用 clock()函数）实现。时钟值首先存储在寄存器中，待核函数执行结束后再写入全局内存，以避免慢速的全局内存访问干扰计时测量结果。

在研究缓存层次结构时，研究人员观察到：穿越互连网络的内存请求，例如访问 L3 缓存和片外内存的请求，其延迟会因执行代码的 TPC 不同而变化。因此，研究人员会对所有 10 个 TPC 位置的测量结果取平均值，并在相关情况下报告延迟的变化范围。

B. 从延迟图推断缓存特征

大多数缓存和 TLB 参数的测量采用步长访问不同大小数组的方式，并绘制平均访问延迟图。本节描述的基本技术也可用于测量 CPU 缓存参数，研究人员针对指令缓存和共享缓存层次结构开发了相应的变体技术。

图 4：3 路 12 行组相联缓存及其延迟图:（a）384 字节、3 路、4 组、32 字节行缓存的延迟图（b）大小为 480 字节（15 行）的数组在缓存中的映射情况

图 4 展示了从平均延迟图中提取缓存大小、路大小（way size）和行大小（line size）的示例。

该示例假设采用最近最少使用替换策略、组相联缓存，且无预取机制。

这些可通过图 4(a)的示例延迟图，来推断缓存参数，具体如下：
* 只要数组大小在缓存容量范围内，延迟就保持恒定（数组大小为 384 字节及以下时）。
* 一旦数组大小超过缓存容量，延迟会逐步增加，增加的步数等于缓存组的数量（4 个组），这是因为缓存组会逐个溢出（数组大小为 385-512 字节时）。
* 触发每步延迟增加所需的数组大小增量等于缓存行大小（32 字节）。
* 当所有缓存组均溢出后（数组大小 ≥16 个缓存行），延迟进入平稳阶段。
* 缓存相联度（3 路）可通过缓存大小（384 字节）除以路大小（128 字节）计算得出。该计算无需知道行大小或缓存组数量。

此外，还存在其他计算这四个缓存参数的方法，因为已知其中任意三个参数，即可通过公式求出第四个参数。

下面代码清单 1 和清单 2 展示了内存微基准测试程序的结构。

    代码清单 1：数组初始化（CPU 代码），  
    for (i = 0; i < array_size; i++) {  
        int t = i + stride;  
        if (t >= array_size) t %= stride;  
        host_array[i] = (int)device_array + 4*t;  
    }  
    cudaMemcpy(device_array, host_array, ...); // 将主机数组数据拷贝到设备数组  

    代码清单 2：依赖读取序列（GPU Kernel代码）  
    int *j = &device_array[0];  
    // 开始计时（start timing）  
    repeat256(j = *(int **)j;); // 宏定义，重复执行256次读取操作  
    // 结束计时（end timing）  

对于每个数组大小和步长，微基准测试程序会执行一系列依赖读取操作（dependent reads），预计算的步长访问模式存储在数组中，从而消除计时内循环中的地址计算开销。
* 步长应小于缓存行大小，以确保能观察到延迟图中的所有步骤；
* 同时步长也应足够大，以确保延迟步骤之间的过渡清晰可辨。

四、测试与结果

本节将详细介绍我们的测试内容及结果。
* 首先，我们测量了clock()函数的延迟；
* 随后研究了流多处理器的各类运算流水线、分支发散与屏障同步机制；
* 此外，还探究了流多处理器内部及周边的内存缓存层级结构，以及内存地址转换与转换后备缓冲器（TLB）。

A. 时钟开销与特性

所有时序测量均使用 clock() 函数，该函数返回一个计数器的值，该计数器每个时钟周期递增一次。clock() 函数会被翻译为“从时钟寄存器中取值，随后执行一次依赖的左移 1 位操作”，这表明该计数器的递增频率是着色器时钟频率的一半。在 clock() 函数之后执行一条非依赖操作，总共需要 28 个时钟周期。

图5：10个和30个线程块的两次连续内核启动时序图。内核调用是串行的，表明线程处理集群（TPC）具有独立的时钟寄存器

图 5 中的实验表明，时钟寄存器是每个线程处理集群（TPC）独有的。图中的数据点表示线程块执行开始和结束时调用 clock() 函数返回的时间戳值。我们观察到：
* 在同一个 TPC 上运行的线程块共享时间戳值，因此也共享时钟寄存器。
* 若时钟寄存器是全局同步的，那么一个内核中所有线程块的启动时间应大致相同。
* 反之，若时钟寄存器是每个流多处理器（SM）独有的，那么同一个 TPC 内的线程块启动时间不会共享相同的时间戳。

B. 运算流水线

图 1：每个流式多处理器包含 8 个标量处理器

每个流多处理器（SM）包含三种不同类型的执行单元，如图 1 和表 I 所示：

表 I：英伟达公开的 GT200 参数

• 8 个标量处理器（SP，Scalar Processor）：执行单精度浮点运算、整数算术运算及逻辑运算。
• 2 个特殊功能单元（SFU，Special Function Unit）：负责执行超越函数（如反平方根、正弦、余弦）和数学函数，同时也处理单精度浮点乘法。
• 1 个双精度单元（DPU，Double Precision Unit）：处理 64 位浮点数的运算。

表II：运算流水线的延迟与吞吐量

表 II 展示了当所有操作数均位于寄存器中时，这些执行单元的延迟与吞吐量。

表III：算术与逻辑运算的延迟与吞吐量

表 III 显示，单精度和双精度浮点数的乘法与乘加运算（mad）均映射为一条设备指令。
* 然而，32 位整数乘法会转换为 4 条原生指令，需要 96 个时钟周期。
* 32 位整数乘加运算则转换为 5 条依赖指令，耗时 120 个时钟周期。硬件仅通过 mul24() 内联函数支持 24 位整数乘法。
* 对于 32 位整数和双精度操作数，除法会转换为子程序调用，导致延迟高、吞吐量低。
* 而单精度浮点数除法会转换为一段简短的内联指令序列，延迟显著更低。

表IV：数学内联函数的延迟与吞吐量（表中“-”表示该操作映射到多指令例程）

单精度浮点数乘法的实测吞吐量约为 11.2 操作数/时钟周期，这一数值高于标量处理器（SP）的 8 操作数/时钟周期，表明乘法操作会同时分发到标量处理器（SP）和特殊功能单元（SFU）。这意味着每个特殊功能单元（SFU）每秒可执行约 2 次乘法（2 个 SFU 总计 4 次），是其他映射到 SFU 的复杂指令吞吐量的两倍。单精度浮点数乘加运算（mad）的吞吐量为 7.9 操作数/时钟周期，这表明乘加运算无法由特殊功能单元（SFU）执行。

通过 decuda（Nvidia 机器级指令的反汇编工具）观察发现，sinf()、cosf() 和 exp2f() 内联函数均转换为两条针对单个操作数的依赖指令序列。CUDA 编程指南指出特殊功能单元（SFU）执行超越函数，但这些超越函数的延迟与吞吐量测量结果，与 SFU 执行的简单指令（如 log2f()）并不匹配。sqrt()（平方根）映射为两条指令：一条反平方根指令和一条倒数指令。

图6：标量处理器（SP）的吞吐量与延迟。6个或7个线程束无法充分利用流水线

图 6 展示了随着流多处理器（SM）上线程束数量的增加，依赖标量处理器（SP）指令（整数加法）的延迟与吞吐量变化。
当并发线程束数量少于 6 个时，实测延迟为 24 个时钟周期。由于所有线程束的延迟相同，表明线程束调度器是公平的。
* 在流水线未饱和时，吞吐量随线程束数量增加呈线性增长。
* 一旦流水线饱和，吞吐量便稳定在 8 操作数/时钟周期（即标量处理器的数量）。

CUDA 编程指南指出，6 个线程束（192 个线程）足以隐藏寄存器的写后读（read-after-write）延迟，但实测发现，当流多处理器（SM）中有 6 个或 7 个线程束时，调度器仍无法填满流水线。

C. 控制流：warp内线程的发散与收敛

1）分支发散

一个线程束中的所有线程在同一时间执行同一条公共指令。CUDA 编程指南指出，当线程束中的线程因数据依赖的条件分支而出现发散时，线程束会串行执行每个被选中的分支路径，并禁用不在该路径上的线程。我们的观察结果与这一预期行为一致。

图7：两个32路发散线程束的执行时序图。上方曲线代表线程束0（Warp0）的时序，下方曲线代表线程束1（Warp1）的时序

图 7 展示了一个线程块中两个并发线程束的实测执行时序，这两个线程束的线程均出现 32 路发散（即每个线程因线程 ID 不同而选择不同路径），且每个线程会执行一段算术运算序列。该图表明，在单个线程束内，每条路径会被串行执行；而不同线程束的执行则可能存在重叠。在一个线程束内，选择同一条路径的线程会并发执行。

2）收敛

当发散路径的执行完成后，线程会收敛到同一条执行路径。通过 decuda 观察发现，编译器会在可能发生发散的分支之前插入一条指令，该指令会向硬件提供收敛点的位置；同时，收敛点处的指令会通过指令编码中的一个字段进行标记。我们观察到，当线程发生发散时，每条路径的执行会串行进行，直到收敛点；只有当一条路径到达收敛点后，另一条路径才会开始执行。

图8：清单3所示内核的执行时序图。数组c包含递增序列{0, 1, …, 31}

根据 Lindholm 等人的研究，GPU 使用分支同步栈（branch synchronization stack）来管理发生发散与收敛的独立线程。

清单 3：收敛栈测试代码
if (tid == c[0]) { ... }
else if (tid == c[1]) { ... }
else if (tid == c[2]) { ... }
...
else if (tid == c[31]) { ... }

我们通过清单 3 所示的内核验证了这一说法。数组 c 包含 0 到 31 之间数字的一个排列，用于指定线程的执行顺序。我们观察到，当线程束遇到条件分支时，“被选中的路径”（taken path）总会优先执行，以数组 c 为递增序列 {0,1,...,31} 为例：
* 最后一个 else if (tid == c[31])（对应 tid=31 的分支）属于“最内层 else 路径关联的 then 子句”，会最先执行。
* 第一个 if (tid == c[0])（对应 tid=0 的分支）属于“最外层 if 的 then 子句”，会最后执行。
* 中间的 else if 分支（如 tid=1 到 tid=30）则按“从内层到外层”的顺序依次执行。

这一现象本质是 GPU 分支同步栈的“落空路径入栈、被选中路径优先执行”机制：每次遇到 if-else，硬件会先执行 else 关联的路径，同时将 if 的 then 路径压入栈，待 else 路径执行完后，再从栈中弹出未执行的 then 路径依次执行，最终实现所有线程的收敛。

图8：清单3所示内核的执行时序图。数组c包含递增序列{0, 1, …, 31}

图 8 展示了当数组 c 为递增序列 {0, 1, ..., 31} 时，该内核的执行时序——此时线程 31 最先执行。当数组 c 为递减序列 {31, 30, ..., 0} 时，线程 0 最先执行，这表明线程 ID 不会影响执行顺序。实测的执行顺序与“被选中的路径优先执行、落空路径（fall-through path）入栈”的机制一致。其他测试表明，一条路径上的活跃线程数量也不会影响路径的执行优先级。

3）SIMT 架构下串行化的影响

CUDA 编程指南指出，为保证正确性，开发者可忽略 SIMT（单指令多线程）行为。本节将展示一个示例：若线程是完全独立的，该代码可正常工作，但在 SIMT 模型下会因线程束行为而死锁。

清单 4：因 SIMT 行为而失效的示例代码
int __shared__ sharedvar = 0;
while (sharedvar != tid) ;
sharedvar++; /* ** 收敛点 ** */

在清单 4 中，若线程独立，第一个线程会跳出 while 循环并递增 sharedvar，随后每个后续线程会依次跳出循环、递增 sharedvar，从而允许下一个线程执行。

但在 SIMT 模型中，当线程 0 不满足 while 循环条件时，会发生分支发散，编译器会将收敛点标记在 sharedvar++ 之前。当线程 0 到达收敛点时，其他被串行化的路径会开始执行；而线程 0 必须等待其余所有线程也到达收敛点后，才能继续执行 sharedvar++——但这些线程永远无法到达收敛点，最终导致死锁。

D. 屏障同步

线程块内不同线程束之间的同步通过 syncthreads() 实现，该函数本质上是一个屏障（barrier），确保所有参与同步的线程都到达该点后才继续执行。对于单个线程束执行的 syncthreads() 序列，该函数被实现为一条单指令，延迟为 20 个时钟周期。

CUDA 编程指南建议，仅当条件在整个线程块内的计算结果完全相同时，才可在条件代码中使用 syncthreads()。本节剩余部分将探究违反该建议时 syncthreads() 的行为：
* 我们将证明 syncthreads() 是以线程束为粒度进行同步的，而非线程。
* 同时将展示，当线程束因分支发散而被串行化时，某一条路径上的 syncthreads() 不会等待另一条路径上的线程，而只会等待同一线程块内运行的其他线程束。

1）单个线程束内线程的 syncthreads()

CUDA 编程指南指出，syncthreads() 是同一线程块内所有线程的屏障。

清单 5：表明`syncthreads()`以线程束为粒度进行同步的示例代码
if (tid < 16) {
    shared_array[tid] = tid;
    __syncthreads(); // 生产者线程的屏障
} else {
    __syncthreads(); // 消费者线程的屏障
    output[tid] = shared_array[tid % 16];
}

但清单 5 中的测试表明，syncthreads() 实际是同一线程块内所有线程束的屏障。该内核由单个线程束执行，其中线程束的前半部分（线程）在共享内存中生成数据，供后半部分（线程）读取。

若 syncthreads() 会等待线程块内的所有线程，那么示例中的两个 syncthreads() 会成为一个公共屏障，迫使生产者线程（线程束前半部分）先写入数据，再由消费者线程（线程束后半部分）读取。

此外，由于分支发散会导致发散线程束的执行串行化（见第 IV-C1 节），若在发散线程束内使用 syncthreads()，内核会发生死锁（在本示例中，16 个线程会等待另 16 个被串行化的线程到达其 syncthreads() 调用点）。但我们的实测结果显示，内核并未死锁，且线程束后半部分并未读取到共享数组（shared_array）中的更新值——这是因为 else 子句会优先执行（见第 IV-C1 节）。这表明，syncthreads() 并不会像 CUDA 编程指南描述的那样，对单个线程束内的发散线程进行同步。

2）多个线程束间的 syncthreads()

syncthreads() 是一个屏障，它会等待所有线程束要么调用 syncthreads()，要么终止。若存在一个既不调用 syncthreads() 也不终止的线程束，syncthreads() 会无限等待——这表明其缺乏超时机制。

清单 6：因`syncthreads()`导致死锁的示例代码（测试使用两个线程束）
// 两个线程束的测试
int count = 0;
if (warp0) { // 线程束0的分支
    count = 1;
    __syncthreads(); // 线程束0的屏障
} else { // 线程束1的分支
    while (count == 0); // 等待count被设为1
    // 线程束1未调用__syncthreads()
}

清单 6 展示了一个无分支发散情况下的死锁示例：第二个线程束会一直循环等待第一个线程束在`syncthreads()`之后生成的数据。

清单 7 详细说明了`syncthreads()`与分支发散之间的相互作用。鉴于`syncthreads()`以线程束为粒度工作，人们可能会预期：要么硬件会忽略发散线程束内的`syncthreads()`，要么**发散线程束会以与非发散线程束相同的方式参与屏障同步**。我们的测试表明，后者是正确的。

```c
清单 7：因`syncthreads()`导致非预期结果的示例代码
if (warp0) { // 线程束0的分支
    // 两路发散
    if (tid < 16) {
        __syncthreads(); // [1] 线程束0内分支1的屏障
    } else {
        __syncthreads(); // [2] 线程束0内分支2的屏障
    }
}
if (warp1) { // 线程束1的分支
    __syncthreads(); // [3] 线程束1的屏障1
    __syncthreads(); // [4] 线程束1的屏障2
}

在该示例中，第二个syncthreads()（[2]）会与第三个syncthreads()（[3]）同步，第一个syncthreads()（[1]）会与第四个syncthreads()（[4]）同步——对于线程束 0，代码块[2]会优先于代码块[1]执行，因为代码块[2]是分支的“被选中路径”（见第 IV-C1 节）。这一结果证实：syncthreads()以线程束为粒度工作，且发散线程束也不例外；每条被串行化的路径会单独执行syncthreads()（代码块[2]不会在屏障处等待代码块[1]），它只会等待线程块内所有其他线程束也执行syncthreads()或终止。

E. 寄存器文件

我们验证了 CUDA 编程指南[1]中的描述：寄存器文件包含 16384 个 32 位寄存器（总计 64 KB）。一个线程使用的寄存器数量会向上取整为 4 的倍数[4]。若尝试启动“每个线程使用超过 128 个寄存器”或“一个线程块使用的寄存器总计超过 64 KB”的内核，启动会失败。

图9：一个线程块使用的寄存器总数限制为16384个（64 KB）。当受寄存器文件容量限制时，线程块内的最大线程数会被量化为64的倍数

在图 9 中：
* 当每个线程使用的寄存器数量少于 32 个时，寄存器文件无法被充分利用——因为每个线程块允许的最大线程数为 512；
* 当每个线程使用的寄存器数量超过 32 个时，寄存器文件容量会限制线程块内可运行的线程数量。

图 9 还表明，当受寄存器文件容量限制时，线程块内的最大线程数会被量化为 64 的倍数。这对内核可使用的寄存器数量施加了额外限制，最明显的案例是“每个线程使用 88 个寄存器”：此时一个线程块仅能运行 128 个线程，仅能使用 11264 个寄存器（128 个线程 ×88 个寄存器/线程），寄存器文件的利用率仅为 69%。

线程块内线程数被量化为 64 的倍数，这表明每个线程的寄存器会分配到 64 个逻辑“存储体”（bank，用于并行访问内存/寄存器以提高带宽）中的一个。每个存储体的大小相同，因此每个存储体可容纳的线程数相同——当受寄存器文件容量限制时，线程数会被限制为 64 的倍数。需注意，这与“寄存器总使用量被量化”的机制不同。

由于 8 个标量处理器（SP）在任意时刻始终执行同一条指令，64 个逻辑存储体的物理实现可在标量处理器之间共享地址线，并使用更宽的存储阵列，而非 64 个独立的物理存储体。每个标量处理器每时钟周期可进行 4 次寄存器访问（对应 4 个逻辑存储体），这一带宽足以支持每时钟周期执行“三读一写”操作数的指令（如乘加指令）。由于一个线程的所有寄存器都位于单个存储体中，该线程需在多个周期内访问其寄存器，而多个线程的寄存器访问可同时进行。

每个标量处理器配备 8 个逻辑存储体，可为“使用特殊功能单元（SFU）的双发射特性”（见第 IV-B 节）以及“算术运算与内存操作并行执行”提供额外带宽。

CUDA 编程指南间接提到“线程数最好是 64 的倍数”：为避免存储体冲突（bank conflict，多个线程同时访问同一存储体导致的访问延迟增加），“最佳结果”需满足“线程块内的线程数是 64 的倍数”。我们观察到，当受寄存器数量限制时，线程块内的线程数会被限制为 64 的倍数，且未观察到存储体冲突。

F. 共享内存

共享内存是每个流多处理器（SM）独有的非缓存内存空间，供同一线程块内的线程共享数据以实现协作。一个线程块允许使用的共享内存大小为 16 KB；此外，内核的函数参数也会占用共享内存空间，因此实际可用的共享内存大小会略有减少。

我们使用清单 1 和清单 2 中的步长访问（stride access，按固定间隔访问数组元素）方法，测得共享内存的读取延迟为 38 个时钟周期。Volkov 和 Demmel 在 GT200 的前代产品 8800GTX 上也测得类似的延迟（36 个时钟周期）[7]。CUDA 编程指南指出，共享内存的延迟与寄存器访问延迟相当。通过改变微基准测试中内存访问的足迹（footprint，访问的内存范围）和步长，我们验证了共享内存确实无缓存。

G. 全局内存

全局内存可被所有运行中的线程访问，即使这些线程属于不同的线程块。全局内存访问是无缓存的，且 CUDA 编程指南文档化的延迟为 400-600 个时钟周期[1]。我们的微基准测试通过“指针追逐式的依赖读取”（类似清单 1 和清单 2 的方式）访问全局内存，在无 TLB（转换后备缓冲器）缺失的情况下，测得读取延迟范围为 436-443 个时钟周期。第 IV-I2 节将详细介绍内存地址转换对全局内存访问延迟的影响；此外，我们还探究了全局内存是否存在缓存，未观察到缓存效应。

H. 纹理内存

纹理内存是一种只读、全局可见的缓存内存空间。在图形渲染中，纹理通常是二维的，且具有二维局部性（spatial locality，相邻地址的数据被连续访问的特性）；CUDA 支持一维、二维和三维纹理。我们以“绑定到线性内存区域的一维纹理”为对象，测量其缓存层级。

图10：纹理内存的缓存延迟。L1缓存为5 KB，L2缓存为256 KB（8路相联）。测试使用64字节步长

测试代码通过依赖纹理读取（类似清单 1 和清单 2 的方式）访问纹理，图 10 展示了纹理内存存在两级缓存：L1 缓存大小为 5 KB，L2 缓存大小为 256 KB（测试使用 64 字节步长）。

我们认为，更高维度（二维、三维）纹理的内存层级不会有显著差异。二维空间局部性通常通过“地址计算将纹理元素重新排列为‘瓦片’（tile，二维数据块）”来实现，而非依赖专用缓存[8]-[10]。

1）纹理 L1 缓存

纹理 L1 缓存为 5 KB，采用 20 路组相联（20-way set-associative，每个缓存组有 20 个缓存行）结构，缓存行大小为 32 字节。

图11：纹理L1缓存的延迟。L1缓存为5 KB（20路相联，32字节缓存行）。测试使用8字节步长。图中还展示了所有线程处理集群（TPC）位置的最大与最小平均延迟：L2缓存的延迟与TPC位置相关

图 11 聚焦于 5 KB 处的第一次延迟跃升，展示了 8 字节步长下的测试结果：对于 5 KB 的缓存，256 字节的“路大小”（way size，每个组中单个路的容量）意味着 20 路组相联。实测显示，L1 缓存命中延迟（261 个时钟周期）是主内存延迟（499 个时钟周期）的一半以上，这与 CUDA 编程指南的描述一致——纹理缓存的作用是减少 DRAM（动态随机存取存储器）带宽需求，而非降低读取延迟。

2）纹理 L2 缓存

纹理 L2 缓存为 256 KB，采用 8 路组相联结构，缓存行大小为 256 字节。

图10：纹理内存的缓存延迟。L1缓存为5 KB，L2缓存为256 KB（8路相联）。测试使用64字节步长

图11：纹理L1缓存的延迟。L1缓存为5 KB（20路相联，32字节缓存行）。测试使用8字节步长。图中还展示了所有线程处理集群（TPC）位置的最大与最小平均延迟：L2缓存的延迟与TPC位置相关

图 10 显示，对于 256 KB 的缓存，32 KB 的“路大小”意味着 8 路组相联。

图12：纹理L2缓存的延迟。L2缓存为256 KB（8路相联，256字节缓存行）。测试使用64字节步长

图 12 放大了图 10 中 256 KB 附近的区域，可观察到“延迟阶梯”——这表明缓存行大小为 256 字节。此外，图 11 还显示，纹理 L2 缓存的访问时间与线程处理集群（TPC）的位置相关，这暗示纹理 L2 缓存并不位于 TPC 内部。

I. 内存地址转换

我们采用类似清单 1 和清单 2 的“步长访问式依赖读取”方法，探究转换后备缓冲器（TLB）的存在；TLB 参数的测量与缓存参数测量类似，但需使用更大的数组大小和与页面大小相当的步长。第 IV-I1 节和第 IV-I2 节分别详细介绍全局内存和纹理内存的 TLB 测试结果。

1）全局内存地址转换

图13：全局内存的TLB延迟。L1 TLB为8 MB全相联，L2 TLB为32 MB 8路相联。测试使用512 KB步长

图 13 显示，全局内存存在两级 TLB：
* L1 TLB 为全相联（fully-associative，所有缓存行可映射到任意缓存组）结构，可缓存 8 MB 内存的映射关系，包含 16 个 TLB 行，每个 TLB 行大小为 512 KB；
* L2 TLB 为 32 MB，采用 8 路组相联结构，TLB 行大小为 4 KB。

本文中，“TLB 大小”指 TLB 可缓存映射关系的内存总大小，而非 TLB 条目（entry）的原始存储大小。例如，“8 MB TLB”表示当页面大小为 4 KB 时，该 TLB 可缓存 2048 个（8 MB / 4 KB）映射关系；若 TLB 行大小为 512 KB，则该 TLB 会组织为 16 个 TLB 行，每个 TLB 行包含 128 个连续页面的映射关系。

在图 13 中，第一个延迟平台（约 440 个时钟周期）对应 L1 TLB 命中（与第 IV-G 节测得的全局内存读取延迟一致）；第二个延迟平台（约 487 个时钟周期）对应 L2 TLB 命中；L2 TLB 缺失时的延迟约为 698 个时钟周期。我们通过“固定元素数量、改变步长”的方式，测得 L1 TLB 的 16 路相联特性。

图14：全局内存L1 TLB的特性。L1 TLB为16路全相联，TLB行大小为512 KB。图中展示了访问16个和17个数组元素时的延迟变化

图 14 展示了访问 16 个和 17 个数组元素时的结果：
* 对于大步长（所有元素映射到同一个缓存组，如 8 MB 步长），访问 16 个元素始终会命中 L1 TLB；
* 而访问 17 个元素时，从 512 KB 步长开始会出现 L1 TLB 缺失。

此外，我们还观察到 L1 TLB 仅包含一个缓存组，这意味着它是“TLB 行大小为 512 KB 的全相联结构”——若存在至少两个缓存组，那么当步长不是 2 的幂且大于 512 KB（缓存路大小）时，部分元素会映射到不同的缓存组；此时访问 17 个元素时，应存在某个步长不会导致 L1 TLB 缺失，但我们在步长超过 512 KB（如 608 KB、724 KB、821 KB）时，从未观察到 L1 TLB 命中。不过，从步长超过 4 MB 的结果可看出，L2 TLB 并非全相联结构。

图13：全局内存的TLB延迟。L1 TLB为8 MB全相联，L2 TLB为32 MB 8路相联。测试使用512 KB步长

图 13 显示，L2 TLB 的“路大小”为 4 MB（从 32 MB 到 36 MB 的延迟变化，见第 III-B 节）；结合 L2 TLB 总大小为 32 MB，可推算出 L2 TLB 的相联度为 8。我们通过扩展测试未发现多级分页（multi-level paging，将虚拟地址分为多个层级进行转换的机制）的证据。

尽管 L1 TLB 的行大小为 512 KB，但 L2 TLB 的行大小仅为 4 KB。我们设计了一个微基准测试：使用两组共 20 个元素（每组 10 个），组内元素的步长为 2 MB，两组元素之间的偏移量为“2 MB + offset”；其中第 i 个元素的地址为：若 i < 10，则地址为 i × 2 MB；否则为 i × 2 MB + offset。为避免 16 路 L1 TLB 掩盖访问延迟，测试需访问超过 16 个元素。由于 L2 TLB 的路大小为 4 MB 且测试使用 2 MB 步长，当 offset 为 0 时，20 个元素会映射到两个 L2 缓存组。

图15：全局内存L2 TLB的特性。L2 TLB行大小为4 KB

图 15 显示，L2 TLB 的行大小为 4 KB。

• 当 offset 为 0 时，20 个元素占用两个缓存组（每组 10 个元素），会导致 8 路组相联 L2 TLB 的冲突缺失。
• 当 offset 超过 4 KB（L2 TLB 行大小）时，每组仅包含 5 个元素，不再导致 L2 TLB 冲突缺失。

尽管页面大小可能小于 L2 TLB 的行大小（4 KB），但我们认为 4 KB 是合理的页面大小选择。需注意，Intel x86 架构主要使用 4 KB 页面的多级分页，而 Intel 系列 GPU 则使用 4 KB 页面的单级分页[9], [11]。

2）纹理内存地址转换

我们采用与第 IV-I1 节相同的方法测量纹理内存 TLB 的配置参数（为简洁起见，此处不再重复方法细节）。

图 16（使用 256 KB 步长）显示，纹理内存存在两级 TLB，可缓存的映射关系大小分别为 8 MB 和 16 MB。其中，L1 TLB 为全相联结构，每个 TLB 行可缓存 512 KB 内存的映射关系；L2 TLB 为 8 路组相联结构，TLB 行大小为 4 KB。在 512 KB 步长下，“按虚拟地址索引的 20 路纹理 L1 缓存”会掩盖纹理 L1 TLB 的特性；此时测得的访问延迟为：TLB 命中（497 个时钟周期）、L1 TLB 缺失（544 个时钟周期）、L2 TLB 缺失（753 个时钟周期）。

J. 常量内存

常量内存包含两个段：一个是用户可访问的段，另一个是编译器生成的常量（如分支条件比较值）使用的段[4]。用户可访问段的大小限制为 64 KB。

图 17 的曲线显示，常量内存存在三级缓存，大小分别为 2 KB、8 KB 和 32 KB。

实测延迟包含两条算术指令（一条地址计算指令和一条加载指令）的延迟，因此原始内存访问时间需减去约 48 个时钟周期——即 L1 缓存命中（8 个时钟周期）、L2 缓存命中（81 个时钟周期）、L3 缓存命中（220 个时钟周期）、L3 缓存缺失（476 个时钟周期）。

我们的微基准测试通过依赖常量内存读取（类似前文代码清单 1 和清单 2 的方式）进行测量。

1）常量 L1 缓存

每个流多处理器（SM）内均包含一个 2 KB 的常量 L1 缓存（见第 IV-J4 节），该缓存采用 4 路组相联结构，包含 8 个缓存组，缓存行大小为 64 字节。“路大小”为 512 字节（2 KB 缓存 / 4 路），进一步证实了 4 路组相联结构。 图 18 展示了这些参数。

2）常量 L2 缓存

每个线程处理集群（TPC）内包含一个 8 KB 的常量 L2 缓存（见第 IV-J4 节和第 IV-J5 节），该缓存与指令内存共享，采用 4 路组相联结构（包含 8 个缓存组），缓存行大小为 256 字节。图 17 中 8192 字节（8 KB）附近的区域展示了这些参数：8 KB 缓存的“路大小”为 2 KB，对应 4 路相联度。

3）常量 L3 缓存

我们观察到一个 32 KB 的全局常量 L3 缓存，供所有线程处理集群（TPC）共享。该缓存采用 8 路组相联结构（包含 16 个缓存组），缓存行大小为 256 字节。

图 17 中 32 KB 附近的区域展示了这些参数。此外，L3 缓存的访问延迟（图 17 中 8-32 KB 的区域）因执行测试代码的 TPC 位置不同而存在显著差异，这暗示 L3 缓存位于“连接 TPC 与 L3 缓存及内存的非均匀互连网络”（non-uniform interconnect，不同节点间访问延迟不同的网络）上。即使数组大小超过 32 KB（需访问主内存），延迟差异也不会随数组大小增加而变化，这 表明 L3 缓存靠近主内存控制器。

我们还测量了 L3 缓存的带宽。 图 19 展示了“不同数量的并发线程块发起 L3 缓存读取请求”时的聚合 L3 缓存读取带宽（每个线程内的请求相互独立）。当线程块数量在 10-20 个之间时，测得的 L3 常量缓存聚合带宽约为 9.75 字节/时钟周期。

我们设计了两种带宽测试变体：
* 一种是每个线程块使用 1 个线程
* 另一种是每个线程块使用 8 个线程（以增加单个 TPC 内的常量缓存读取需求）。

两种测试在线程块数量少于 20 个时表现出相似的带宽特性，这表明：
* 当仅运行 1 个线程块时，即使增加线程块内的读取需求（通过多线程），一个流多处理器的读取带宽也仅能达到约 1.2 字节/时钟周期。
* 在线程块数量超过 20 个的“8 线程/线程块”场景中，测试结果无效——因为此时缺乏足够的唯一数据集，且每个 TPC 的 L2 缓存会掩盖部分 L3 缓存请求，导致表观聚合带宽增加；
* 当线程块数量超过 30 个时，部分流多处理器会运行多个线程块，导致负载不均衡。

4）缓存共享机制

常量内存的 L1 缓存为每个流多处理器（SM）独有，L2 缓存由同一线程处理集群（TPC）内的所有 SM 共享，L3 缓存为全局共享。

我们通过“测量两个并发线程块在不同位置的延迟（同一 SM、同一 TPC、不同 TPC）”验证了这一机制： 两个线程块会竞争共享缓存，导致实测缓存大小减半。

图 20 展示了测试结果：
* 在所有场景中，全局 L3 缓存的实测大小均减半至 16 KB；
* 当两个线程块位于同一 TPC 时，TPC 级 L2 缓存的实测大小减半至 4 KB；
* 当两个线程块位于同一 SM 时，SM 级 L1 缓存的实测大小减半至 2 KB。

5）与指令内存的缓存共享

已有研究提出，常量缓存与指令缓存的部分层级是统一的[12], [13]。我们的测试证实： L2 和 L3 缓存是“指令-常量统一缓存”，而 L1 缓存是专用缓存（仅用于常量或仅用于指令）。

我们采用与第 IV-J4 节类似的方法，通过“测量不同位置的线程块在‘指令读取’与‘常量缓存读取’之间的干扰” 验证了这一点，结果 如图 21 所示：即使线程块在同一 SM 上运行，指令读取需求也不会影响 L1 缓存的访问时间，这表明 L1 缓存是专用的。

K. 指令供给

我们检测到指令缓存存在三级结构，大小分别为 4 KB、8 KB 和 32 KB，如图 22 所示。

分析指令缓存层级，关键有 3 点：
1. 测试代码设计逻辑：用“不同大小的独立 8 字节 abs 指令块”构建，核心目的是“最大化指令读取需求”——通过独立指令块避免指令复用，强制 GPU 持续从缓存/内存读取新指令，从而暴露底层缓存的真实容量与层级特征，排除“指令重复使用掩盖缓存行为”的干扰。
2. 实测缓存层级结果：从图 22 中明确观察到两级指令缓存：
* L2 缓存：容量为 8 KB，是指令存储层级的中间环节；
* L3 缓存：容量为 32 KB，作为更全局的指令缓存层。
但未检测到理论中可能存在的 4 KB L1 指令缓存（GPU 常见的最接近执行单元的高速指令缓存）。
3. L1缓存未检出的原因：推测是“少量指令预取”的影响，GPU 硬件会提前将即将执行的指令加载到缓存（预取机制），即便存在 4 KB L1 缓存，预取操作也会让指令提前进入 L1，使得“L1 缓存命中延迟”与“预取后的指令访问延迟”难以区分，最终导致测试无法捕捉到 L1 缓存的独立存在特征，掩盖了其层级痕迹。

1）指令 L1 缓存

每个流多处理器（SM）内包含一个 4 KB 的指令 L1 缓存，采用 4 路组相联结构，缓存行大小为 256 字节。

为测量 L1 缓存参数（如图 23 所示），我们在同一 TPC 的另外两个 SM 上运行并发线程块，以引入 L2 缓存竞争——这样，从 4 KB 开始的 L1 缓存缺失便不会像图 22 那样被掩盖。图 23 中可观察到 256 字节的缓存行大小以及 4 个缓存组的存在。指令 L1 缓存为每个 SM 独有：当同一 TPC 内的其他 SM 填满其指令缓存层级时，被观测 SM 的 4 KB 指令缓存大小并未减少。

2）指令 L2 缓存

每个线程处理集群（TPC）内包含一个 8 KB 的指令 L2 缓存，采用 4 路组相联结构，缓存行大小为 256 字节。第 IV-J5 节已证实，该 L2 缓存同时用于指令内存和常量内存；我们通过测试验证了指令 L2 缓存的参数与常量 L2 缓存一致。

3）指令 L3 缓存

指令 L3 缓存为全局共享，大小为 32 KB，采用 8 路组相联结构，缓存行大小为 256 字节。第 IV-J5 节已证实，该 L3 缓存同时用于指令内存和常量内存；我们通过测试验证了指令 L3 缓存的参数与常量 L3 缓存一致。

4）指令读取

流多处理器（SM）似乎会从指令 L1 缓存中以 64 字节为单位读取指令（对应 8-16 条指令）。图 24 展示了“36 次连续 clock() 读取（共 72 条指令）”的执行时序，该结果为 10000 次微基准测试执行的平均值。

当一个线程束运行测量代码时，同一 SM 上运行的 7 个“驱逐线程束”（evicting warps）会通过“循环执行 24 条指令（192 字节）”反复驱逐图中大点所示区域的缓存行——这些指令会导致指令缓存的冲突缺失。
* 驱逐线程束是否会反复驱逐测量代码使用的缓存行，取决于线程束调度器，且概率较高；
* 仅当指令读取跨越边界（图 24 中的 160 字节、224 字节、288 字节、352 字节、416 字节）时，才会观察到驱逐导致的缓存缺失延迟。

我们观察到：当发生冲突时，整个缓存行会被驱逐（覆盖 160-416 字节的代码区域）；且缓存缺失的影响仅在 64 字节的读取块之间体现，在读取块内部无影响。

五、相关工作

微基准测试（microbenchmarking，通过短小精悍的代码测量硬件特定特性的方法）在过去被广泛用于确定各类处理器结构的硬件组织。本节将重点关注针对 GPU 的相关工作。
* Volkov 和 Demmel 对 GT200 的前代产品 8800GTX GPU 进行了基准测试[7]。他们测量了与“加速稠密线性代数运算”相关的 GPU 特性，揭示了纹理缓存的结构和一级 TLB（转换后备缓冲器）的存在。尽管他们的测试基于前代硬件，但其测量结果与我们的结果基本一致。 我们的研究则聚焦于 GPU 的微架构，揭示了额外的 TLB 层级与缓存结构，以及处理核心的组织方式。
* 另有研究通过基准测试对 GPU 进行性能分析，例如 GPUBench[14]——这是一套基于 OpenGL ARB 着色语言（OpenGL ARB shading language，用于编写 GPU 着色器的高级语言）的微基准测试集，可测量 GPU 的部分指令与内存性能特性。
* 然而，相较于 CUDA，更高抽象层级的 ARB 着色语言与硬件的距离更远，难以从测试结果中推断出详细的硬件结构；但 ARB 着色语言具有 CUDA 所不具备的“跨厂商兼容性”优势。

目前，NVIDIA GPU的规格说明和CUDA优化技术主要来自厂商文档[1], [3]。相关的优化研究（如[15]）和性能模拟器（如[2]）均依赖这些公开文档。我们的研究提供了更详细的硬件参数，有望帮助提升这些研究的准确性。

六、总结与结论

本文介绍了我们对 NVIDIA GT200 GPU 的分析及测量方法。我们设计的微基准测试集揭示了 GPU 处理核心与内存层级的架构细节。

虽然 GPU 是一种复杂的设备，我们无法逆向工程出其所有细节，但我们相信已探究了其中具有重要意义的部分特性。表 V 总结了我们的架构发现。

表V：GT200架构总结

我们的结果验证了 CUDA 编程指南[1]中记录的部分硬件特性，同时也发现了一些未文档化的硬件结构 ，例如 控制流机制、缓存与 TLB 层级 。此外，在部分场景中，我们的发现与文档化的特性存在差异 （如纹理缓存和常量缓存）。

我们还介绍了用于架构分析的方法，相信这些方法将对“其他类 GPU 架构的分析 ”及“类 GPU 性能模型的验证 ”有所帮助。

本研究的最终目标是更深入地理解 GPU 硬件，从而充分挖掘其性能潜力。

参考文献

关注“鲸栖”小程序，掌握最新AI资讯