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

关键词：GPU 微架构、微基准测试、CUDA、存储层次、算术流水线、控制流

本文是系列文章《Demystifying GPU Microarchitecture through Microbenchmarking》的第一篇，也是早期 NVIDIA GPU 架构分析文章之一。由于全文篇幅较长（约 2 万字），可能更适合作为参考资料，建议读者根据目录选择感兴趣的部分阅读。

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

更重要的是，本文介绍了一套用于 GPU 架构分析的方法论，相信这些方法对其他类 GPU 架构的分析以及类 GPU 性能模型的验证均会有所助益。

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

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

英伟达 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
全文约 1.7 万字，预计阅读时间 60 分钟，播客音频约 17 分钟。

《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 性能的极致挖掘？

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

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

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

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

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

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

抽象与实现的权衡，而非设计缺陷。论文证实 syncthreads() 以 warp 为同步粒度，但文档描述为“线程级屏障”，本质是 CUDA 编程模型的“高层抽象简化”——将复杂的硬件同步逻辑（warp 调度、分支收敛）封装为统一接口，让开发者无需理解 SIMT 架构的底层细节即可编写代码。这种抽象符合 GPU“海量线程并行”的设计目标，若强制按线程级同步，会大幅增加硬件调度开销，降低并行效率，反而违背性能设计初衷。
修正描述可能引发生态混乱 。论文指出，依赖“线程级同步”假设的代码可能出现死锁，但大量现有 CUDA 程序已基于文档描述开发。若修正描述，需开发者重新适配代码（如调整分支结构、同步时机），不仅增加开发成本，还可能导致部分旧代码失效。厂商选择维持现有描述，是优先保障生态兼容性而非绝对准确性。
引导开发者依赖“规则而非实现”。论文提到厂商通过“模糊规则”（如“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，Translation Lookaside Buffer，用于加速虚拟地址到物理地址转换的硬件缓存）层次结构、常量内存（constant memory）、纹理内存（texture memory）和指令内存缓存。
讨论所采用的测量技术，研究人员认为 这些技术将有助于其他 GPU 及类 GPU 系统的分析与建模 ，并提升 GPU 性能建模与仿真的保真度[2]（fidelity，指模型或仿真结果与真实硬件行为的吻合程度）。

本文其余部分结构如下：

第二节回顾 CUDA 计算模型；
第三节描述测量方法；
第四节呈现测量结果；
第五节回顾相关工作；
第六节总结研究发现。

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

A. GPU 架构

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

这些线程被映射到硬件资源层次结构上。线程块在线式多处理器（SM，Streaming Multiprocessor，GPU 的核心计算单元，负责执行线程块）内执行，如图 1 所示。

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

尽管编程模型中使用的是标量线程集合，但 SM 更接近一个 8 路向量处理器 （ 8-wide vector processor，指一次可处理 8 个数据元素的向量运算单元 ），该处理器对 32 路向量进行操作 。

CUDA 编程模型呈现给开发者的是“操控大量独立标量线程”的抽象，但其底层硬件——流多处理器（SM）——本质上是一个8路向量处理器，能够并行处理8个数据元素（即“8-wide”）。

这意味着，硬件并非为每个标量线程单独分配执行单元，而是以向量运算的方式批量处理数据。CUDA的编程抽象掩盖了这一向量操作的细节，从而降低了开发门槛。

SM内执行的基本单位是线程束（warp）。
* 在GT200架构中，一个线程束由32个线程组成，它们被分为4个8线程的子组（sub-warp），每个子组在8个标量处理器（SP）上并行执行。

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

英伟达将这种架构称为单指令多线程（SIMT）：线程束中的所有线程同步执行相同的指令，但允许每个线程拥有独立的分支路径。
* SM包含算术单元，以及为线程块和线程私有的资源，例如每个线程块专用的共享内存（shared memory）和寄存器文件（register file，用于存储线程执行过程中的临时数据）。

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

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

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

从CUDA的视角看，GPU由TPC集合、互连网络（连接不同硬件单元的通信链路）和存储系统（DRAM内存控制器）构成，如图3所示。表I列出了英伟达公开的GT200关键参数[1, 3]。

表 I：英伟达公开的 GT200 参数[1, 3]

B. CUDA 软件编程接口

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

在CUDA模型中，主机（CPU）代码通过调用设备（GPU）函数来启动GPU核函数。由于GPU与CPU使用不同的指令集，CUDA编译流程会使用不同的编译器分别编译CPU和GPU代码。
* GPU代码首先被编译为PTX中间代码（Parallel Thread Execution，并非最终执行的机器码）。
* 随后，PTX代码被“汇编”为可在GPU硬件上直接执行的原生机器指令。
* 编译后的CPU代码和GPU代码最终被合并为一个单一的“胖二进制文件”（包含适配多种架构的代码）[4]。

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

研究人员发现，最有效的方式是使用CUDA C编写代码，然后通过decuda工具[5]（一种基于对Nvidia编译器输出分析而开发的反汇编工具，用于解析未公开的原生机器指令）在原生代码层面验证生成的机器指令序列。

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

三、测量方法

A. 微基准测试方法

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

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

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

在研究缓存层次结构时，研究人员观察到：穿越互连网络的内存请求（例如访问L3缓存和片外显存的请求），其延迟会因执行代码的TPC不同而变化。因此，研究人员会对所有10个TPC位置的测量结果取平均值，并在相关情况下报告延迟的变化范围。
* 原因分析：这可能是由于互连网络是“非均匀”的。10个TPC在GPU芯片上的物理位置不同，导致它们访问共享的L3缓存或内存控制器的路径长度和延迟存在差异。

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

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

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

图4展示了如何从平均访问延迟图中提取缓存大小、路大小（组相联缓存中每个组包含的缓存块总数）和行大小（缓存与下一级存储之间一次传输的数据块大小）。

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

最近最少使用替换策略，LRU，Least Recently Used，一种缓存替换策略，缓存满时替换最近最少访问的缓存块。
组相联缓存，set-associative cache，将缓存分为多个组，每个组含多个路，内存块映射到特定组并可存入该组任意一路。
预取，prefetching，硬件或软件主动将可能即将访问的数据加载到缓存的技术。

通过图 4(a) 的示例延迟图，可以推断缓存参数，具体如下：

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

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

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

[code]
代码清单 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, …); // 将主机数组数据拷贝到设备数组
[/code]

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

对于每个数组大小和步长，微基准测试程序会执行一系列依赖读取操作（dependent reads，指后一次读取的地址依赖于前一次读取的结果，确保每次读取需等待前一次完成以准确测量延迟），预计算的步长访问模式存储在数组中，从而消除计时内循环中的地址计算开销。

步长应小于缓存行大小，以确保能观察到延迟图中的所有步骤；
同时步长也应足够大，以确保延迟步骤之间的过渡清晰可辨。

四、测试与结果

本节将详细介绍我们的测试内容及结果。

首先，我们测量了 clock() 函数的延迟；
随后研究了流多处理器的各类运算流水线、分支发散与屏障同步机制；
此外，还探究了流多处理器内部及周边的内存缓存层级结构，以及内存地址转换与转换后备缓冲器（TLB，Translation Lookaside Buffer，用于加速虚拟地址到物理地址转换的高速缓存）。

A. 时钟开销与特性

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

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

图 5 中的实验表明，时钟寄存器是每个线程处理集群（TPC）独有的。图中的数据点表示线程块执行开始和结束时调用 clock() 函数返回的时间戳值。我们观察到：

在同一个 TPC 上运行的线程块共享时间戳值，因此也共享时钟寄存器。
若时钟寄存器是全局同步的，那么一个内核中所有线程块的启动时间应大致相同；
反之，若时钟寄存器是每个流多处理器（SM）独有的，那么同一个 TPC 内的线程块启动时间不会共享相同的时间戳。

B. 运算流水线

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

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

表 I：英伟达公开的 GT200 参数[1, 3]

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 操作数/时钟周期（即标量处理器的数量）。

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

1）分支发散

一个线程束中的所有线程在同一时间执行同一条公共指令。

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

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

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

2）收敛

当发散路径的执行完成后，线程会收敛到同一条执行路径。

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

根据 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()，根据 CUDA 编程指南的描述，内核会发生死锁（在本示例中，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()` 导致非预期结果的示例代码
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
}

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

在该示例中，第二个 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字节步长。

测试代码通过依赖纹理读取访问纹理，图 10 展示了纹理内存存在两级缓存：L1 缓存大小为 5 KB，L2 缓存大小为 256 KB。

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

1）纹理 L1 缓存

纹理 L1 缓存为 5 KB，采用 20 路组相联结构，缓存行大小为 32 字节。

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

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

2）纹理 L2 缓存

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

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

图 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. 内存地址转换

我们采用“步长访问式依赖读取”方法，探究转换后备缓冲器（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 为全相联结构，可缓存 8 MB 内存的映射关系，包含 16 个 TLB 行，每个 TLB 行大小为 512 KB。
* L2 TLB 为 32 MB，采用 8 路组相联结构，TLB 行大小为 4 KB。

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

在图 13 中，第一个延迟平台（约 440 个时钟周期）对应 L1 TLB 命中；第二个延迟平台（约 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 的全相联结构”。从步长超过 4 MB 的结果可看出，L2 TLB 并非全相联结构。

图 13 显示，L2 TLB 的“路大小”为 4 MB；结合 L2 TLB 总大小为 32 MB，可推算出 L2 TLB 的相联度为 8。我们通过扩展测试未发现多级分页的证据。

尽管 L1 TLB 的行大小为 512 KB，但 L2 TLB 的行大小仅为 4 KB。我们设计了一个微基准测试：使用两组共 20 个元素（每组 10 个），组内元素的步长为 2 MB，两组元素之间的偏移量为“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：纹理内存的TLB延迟。 L1 TLB为8 MB全相联，L2 TLB为16 MB 8路相联。L1 TLB缺失延迟为544个时钟周期，L2 TLB缺失延迟为753个时钟周期。测试使用256 KB步长。

图 16（使用 256 KB 步长）显示，纹理内存存在两级 TLB，可缓存的映射关系大小分别为 8 MB 和 16 MB。其中，L1 TLB 为 16 路全相联结构，每个 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：常量内存的缓存延迟。 L1缓存为2 KB，L2缓存为8 KB（4路相联），L3缓存为32 KB（8路相联）。测试使用256字节步长。图中还展示了所有线程处理集群（TPC）位置的最大与最小平均延迟：L3缓存的延迟与TPC位置相关。

图 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 展示了这些参数。

图 18：常量L1缓存的延迟。 L1缓存为2 KB（4路相联，64字节缓存行）。测试使用16字节步长。

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 缓存靠近主内存控制器。

图 19：常量L3缓存的带宽。 带宽约为9.75字节/时钟周期。图中展示了“每个线程块1个线程”和“每个线程块8个线程”两种场景的带宽变化。

我们还测量了 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：常量内存的缓存共享机制。 L1缓存为每个SM独有，L2缓存为每个TPC独有，L3缓存为全局共享。测试使用256字节步长。

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

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

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

图 21：常量内存与指令内存的缓存共享机制。 L2和L3缓存与指令内存共享，L1缓存为专用。测试使用256字节步长。

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

K. 指令供给

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

图 22：指令缓存的延迟。 L2缓存为8 KB（4路相联），L3缓存为32 KB（8路相联）。本测试未检测到4 KB的L1缓存。

分析指令缓存层级，关键有 3 点：

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

1）指令 L1 缓存

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

图23：指令L1缓存的延迟。L1缓存为4 KB（4路相联，256字节缓存行）。通过引入L2缓存竞争，使L1缓存缺失可见。图中还展示了所有线程处理集群（TPC）位置的最大与最小平均延迟：L3缓存的延迟与TPC位置相关。

为测量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次微基准测试执行的平均值。

图24：指令读取大小。SM似乎从L1缓存中以64字节为单位读取指令。代码跨越3个256字节的缓存行，缓存行边界位于160字节和416字节处。

当一个线程束运行测量代码时，同一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资讯

本文由鲸栖原创发布，未经许可，请勿转载。转载请注明出处：http://www.itsolotime.com/archives/14084