Uncovering Real GPU NoC Characteristics: Implications on Interconnect Architecture

作者/机构: Zhixian Jin (KAIST), Christopher Rocca (KAIST), Jiho Kim (KAIST), Hans Kasan (KAIST), Minsoo Rhu (KAIST), Ali Bakhoda (Microsoft), Tor M. Aamodt (University of British Columbia), John Kim (KAIST)

A1 主要贡献

本文针对现代NVIDIA GPU中的片上网络（NoC）进行了详尽的延迟和带宽分析，旨在揭示真实GPU互连架构的特性，并探讨其对系统性能和安全性的影响。

核心问题: 现有关于GPU NoC的研究大多基于仿真，其假设可能无法准确反映真实GPU硬件的特性。随着GPU核心（流式多处理器，SMs）数量和内存带宽的持续增长（如表I所示），连接这些组件的NoC已成为决定整体系统性能的关键因素。因此，对真实GPU NoC进行深入分析至关重要。

研究目标:
1. 对多代现代NVIDIA GPU（V100, A100, H100）的NoC进行详细的特性分析，包括延迟和吞吐量（带宽）。
2. 基于分析结果，探讨GPU NoC特性对两个重要方面的影响：时序侧信道攻击和GPU NoC微架构设计。

创新点/主要贡献:
* 对真实GPU NoC的详细特性分析：本文提供了跨越多代NVIDIA GPU的真实片上网络（NoC）的详细特性描述和分析。一个关键发现是，核心（SM）与内存分区之间的延迟呈现非均匀性，而带宽则近似均匀。
* 揭示NoC特性对时序侧信道攻击的影响：基于NoC的特性分析，本文展示了非均匀延迟如何被利用于时序侧信道攻击，同时也揭示了它如何能被用作防御时序侧信道攻击的机制。
* 对GPU架构探索中NoC假设的影响：本文展示了在GPU架构探索中，对NoC做出正确假设的重要性，以确保片上网络（或片上带宽）不会成为整体系统性能的瓶颈。

TABLE I: NVIDIA GPU的微架构比较。

A3 背景知识

A. GPU 架构

GPU的并行处理能力：GPU是一种高度并行的处理器，能够同时执行数千个线程，以实现高吞吐量处理【12】。一个GPU由多个核心或SM（流式多处理器）和内存分区组成，它们通过片上网络（NoC）互连。一个使用如CUDA/OpenCL等GPU特定编程模型编写的GPU上下文包含多个线程块，由线程块调度器【13】进行调度。在GPU流水线中，一个warp（一组线程）是执行的基本单位，它以不同的数据执行相同的指令。此外，多个warp根据warp调度算法【14】–【17】并发执行。

GPU的内存层级：每个核心拥有一个私有的L1缓存和共享内存，而一个全局共享的L2缓存靠近内存控制器。每个内存分区（MP）包含多个L2 slice和一个管理片外内存的内存控制器（MC）【18】，【19】。NoC负责连接SM与内存分区（MP）。

GPU的层级组织：GPU架构是层级化的：在一个纹理处理集群（TPC）内，两个SM通过一个多路复用器连接；7到9个TPC组成一个图形处理集群（GPC）；6到8个GPC组成一个GPU【12】，【20】，【21】。近期的GPU（例如A100, H100）包含一个中央互连，将GPU分割成左、右两个“分区”。本文旨在分析GPU互连，并探讨NoC在延迟和带宽方面对高吞吐量处理器的影响。

B. 相关工作

现有GPU系统分析的局限性：尽管已有对近期GPU系统（包括V100【22】和A100【23】）的分析，以及一些分析GPU内存层级【24】的工作，但这些工作虽然深入剖析了GPU的各个组件，却未对GPU内部的互连进行详细分析。关于多GPU系统的分析【25】–【27】也已存在，旨在理解GPU到GPU的互连影响，但它们对片上互连的讨论有限。

GPGPU的NoC设计挑战：与多核CPU的NoC相比，用于高吞吐量处理器的片上网络（NoC）面临着不同的挑战。针对GPGPU的吞吐量有效NoC【28】被提出，其中基于“多对少对多”（many-to-few-to-many）的流量模式，识别出“回复”网络是瓶颈。为解决此瓶颈，已提出了多种替代性NoC架构【29】–【32】以最小化回复网络的影响。为减少片上流量，提出了合并（coalescing）技术以减少注入NoC的片上流量【33】。然而，这些工作大多假设采用如2D-mesh拓扑结构的多跳网络，这可能导致片上带宽瓶颈。

对现代GPU NoC架构的假设：近期的工作【34】，【35】已经假设了层级交叉开关（hierarchical crossbar）的组织结构。据我们所知，真实系统中的GPU互连与层级交叉开关的组织方式相似；然而，这些工作并未提供对真实GPU NoC的详细分析。更重要的是，许多先前工作没有考虑NoC带宽对L2带宽的影响——即NoC带宽需要被充分配置，以确保能提供足够的L2带宽。Accel-sim【36】也指出，L2带宽可能是一个瓶颈，这阻碍了他们的模拟器在现代GPU上进行更精确的建模。

C. 方法论

实验平台与方法：在这项工作中，我们使用了前面（表I）描述的三款GPU。为了分析GPU片上网络（NoC）的性能影响，我们编写了合成的微基准内核，以隔离GPU内部不同因素和组件的影响。具体来说，测量重点在于理解NoC的延迟和带宽。

1) 延迟测量

延迟测量算法概述：延迟测量的高级概述如算法1所示。我们仅利用一个warp内的一个线程发送请求，以最小化GPU中其他组件的影响——因此，GPU内部不存在合并（coalescing）和任何资源竞争。所做的内存访问保证绕过L1缓存，并且除非另有说明，工作集能完全放入L2缓存。对D[]的内存访问用于预热L2缓存，以确保所有访问都在L2中命中。测量的（往返）延迟包括核心（SM）、NoC和L2访问延迟。由于测量是在不同的SM和L2 slice之间进行的，核心或L2的延迟大致相同，但整体延迟的变化来自于NoC的影响。内存访问通过-dlcm=cg编译选项绕过每个SM的L1缓存。延迟是使用每个SM内部提供的硬件时钟计数器clock()来测量的。

算法1 延迟测量算法。

源与目的地的确定：为了理解NoC的影响，需要知道数据包（或内存请求）的源和目的地。源（即SM）由smid寄存器确定，并用于在必要时将内核“钉”在特定的SM上。目的地（即L2 slice）可以通过使用带有“non-aggregated”模式的nvprof来监控每个slice的流量来确定。该分析器还用于收集访问同一L2 slice的地址信息。从slice ID s到D[]中相应索引idx的映射信息存储在M[]中（即M[s] = idx）。请注意，SM或L2 slice ID的数值本身不重要，只需要确保访问的是不同的SM或L2 slice（对于带宽测量，则是相同的）。

评估范围和观察：初步评估是在V100上完成的，但在A100和H100上也观察到了类似的趋势。然而，更大的GPU（即A100、H100）具有多个“分区”，它们在某些特性上有所不同，我们也将对此进行分析。

2) 带宽测量

带宽测量算法：我们编写了微基准测试来测量片上（或L2 fabric）带宽，如算法2所述，这与常用于测量带宽的流式基准测试非常相似。我们使用了顺序、跨步的内存访问模式，并且与延迟测量类似，绕过了L1缓存。在测量L2 fabric带宽时，L2缓存被预热以确保所有访问都在L2中命中。除非另有说明，执行的是读访问。

带宽饱和与目的地控制：与延迟测量不同，为了使带宽饱和，我们利用了一个warp内的所有线程以及多个线程块——因此，每个线程都执行算法2中描述的内核。算法2的一个独特之处在于，目的地（或L2 slice）的控制方式与延迟测量相似。在评估之前，确定了一组用于数据（D[]）的地址，并向同一个L2 slice（s）发出多次访问以评估L2带宽——因此，M[s][]包含了发送请求到同一L2 slice s的线程所使用的索引。我们进行了不同的L2带宽测量，包括每个L2 slice、每个内存分区和聚合（总）L2带宽。性能（或带宽）是根据传输的数据（消息）大小和执行微基准测试所需的时间来确定的。

算法2 带宽测量算法。

1: function L2 BANDWIDTH(s) ▷ s: target L2 slice
2:     stride ← blockDim.x;
3:     tid ← blockIdx.x × stride + threadIdx.x;
4:     index ← M [s][tid];
5:     ldcg(&D[index]);

A2 方法细节

III. 片上延迟分析

本节我们对GPU的片上（片上网络）延迟进行分析。我们将展示GPU中如何观察到非均匀延迟，这通常由核心和内存分区（MP）（以及L2 slice）的物理位置决定。此外，不断增大的GPU尺寸导致单片GPU包含多个“分区”，我们将分析这类GPU分区对片上延迟的影响。

A. 非均匀延迟

L2访问延迟的非均匀性：图1展示了L2访问延迟，即从SM到L2 slice的往返延迟，包括核心（SM）延迟、NoC延迟和L2访问延迟。L2访问延迟是通过一次在本地L1缓存中未命中但在共享L2缓存中命中的内存读取访问来测量的。图1(a)绘制了V100 GPU上从一个特定核心（SM 24）到32个不同L2 slice的L2访问延迟分布，x轴是NVIDIA分析器【37】提供的L2 slice ID。尽管该图只显示了一个SM，但其他SM也存在非均匀的延迟分布，最低延迟为175个周期，最高延迟为248个周期，平均延迟约为212个周期。对不同L2 slice的访问共享相同的组件（例如，核心延迟，L2延迟），但主要差异在于L2 slice（以及SM位置）的物理位置。


图 1: (a) V100 GPU上从一个SM（即SM 24）到所有L2 slice的非均匀L2访问延迟，以及(b) 每个GPC内的平均延迟和延迟变化。

观察 #1: 从核心（SMs）通过GPU NoC到各个L2 slice的延迟是非均匀的。

GPC间的平均延迟与内部变化：总体而言，V100上的平均L2访问延迟约为212个周期，这与先前工作报告的193个周期的L2延迟【22】相似。如图1(b)所示，虽然各GPC的平均L2延迟是一致的，但每个GPC内部的变化可能很大。例如，GPC2和GPC3的延迟分布较窄，而其他GPC的分布则较宽，因为GPU内的L2延迟对于GPC4来说，差异可达71个周期——这代表了大约33%的L2延迟差异。

GPC延迟分布直方图：图2显示了两个不同GPC的L2延迟直方图，其中收集了V100 GPU上从GPC内所有SM到所有内存L2 slice的延迟数据。虽然两个GPC的平均延迟非常相似，但延迟变化却大不相同——即GPC0（µ = 213个周期，σ = 13.9个周期）和GPC2（µ = 209个周期，σ = 7.5个周期）。因此，L2延迟根据源（SM）和目的地（L2 slice）的不同而变化。


图 2: 两个不同GPC的L2延迟直方图分布，(a) GPC0和(b) GPC2。两个GPC的平均L2延迟非常相似，但延迟分布不同。

观察 #2: 来自一个GPC的平均L2延迟（即从GPC内的SM到所有L2 slice的延迟）在所有GPC之间是相似的；然而，GPC内的SM之间存在L2延迟变化，并且延迟变化在不同GPC之间有所不同。

B. 物理布局与延迟

物理位置对延迟的决定性作用：如第二节所述，GPU是层级化组织的，多个核心位于一个GPC内，多个L2 slice位于一个MP内。L2访问延迟（以及片上网络延迟）的精确值由核心（或SM）在GPC内的物理位置以及L2 slice在MP内的物理位置决定——此外还取决于GPC和MP本身的位置。为了更好地理解布局的影响，我们在图3中绘制了L2延迟，但首先根据它们的MP ID将L2 slice ID分组——即位于同一MP内的L2 slice被分在一起。然后根据延迟值对L2 slice ID进行排序——从最低的L2延迟到最高的。不同SM的slice ID排序顺序是相同的——例如，对于MP 0，所有四个图都按以下L2 slice ID顺序显示L2延迟：0, 4, 27, 31, ...。为便于说明，仅显示了四个SM的延迟分布，其中每组两个SM来自同一个GPC。虽然确切的延迟值不同，但来自同一GPC的SM在延迟变化趋势上表现出一致性（例如，图3(a,b)和图3(c,d)）。虽然不同MP之间的延迟不同，但每个MP内部的延迟趋势在不同SM之间也是相似的。


图 3: 多个SM的延迟变化图，包括来自GPC0的两个SM（(a) SM60 和 (b) SM24）和来自GPC4的两个SM（(c) SM64 和 (d) SM28）。L2 slice索引经过排序，因为同一内存分区中的L2 slice被分组在一起。

V100逻辑布局图分析：基于芯片照片【38】和我们的分析，V100的一个近似逻辑布局图如图4所示。该图包含了一个GPC（即GPC 4）中SM的近似布局和一个内存分区（MP）（即MP 3）中L2 slice的布局。实际的SM ID或L2 slice ID并不重要，但该图显示了核心（SM）和L2 slice的相关布局。物理位置相近的SM和L2 slice延迟最低（例如，L2 slice 17和SM64的L2延迟为180个周期），而距离最远的（例如，L2 slice 15和SM4）延迟最高，为217个周期。为简化起见，该图仅显示了一个GPC和一个MP，但这一趋势在不同的GPC和MP组合中都成立。


图 4: V100 GPU的框图（近似逻辑布局），显示了MP内的L2 slice布局和GPC内的SM布局。

跨SMs和L2 slices的延迟细节：图5显示了不同SMs和L2 slices之间的详细延迟分析。结果表明，不同的SM位置导致L2延迟存在一个恒定的差异，而某些L2 slices总是比其他L2 slices的延迟更低。


图 5: V100上GPC4中的SM到内存分区（MP）3内不同L2 slice的L2访问延迟。物理上更近的SM和L2 slice具有更低的延迟。

观察 #3: 非均匀的L2延迟由SM在GPC内的物理位置和L2 slice在内存分区内的物理位置决定。

皮尔逊相关性定义：皮尔逊相关性【39】是衡量两个数据集之间线性相关性的常用指标，可应用于测量两个SM之间L2延迟分布的相似性。该相关性（r）定义如下：

$$r=\frac{\sum_{i=1}^{n}(X_{i}-\overline{X})(Y_{i}-\overline{Y})}{\sqrt{\sum_{i=1}^{n}(X_{i}-\overline{X})^{2}}\sqrt{\sum_{i=1}^{n}(Y_{i}-\overline{Y})^{2}}}$$

其中n是样本数，Xi, Yi是来自两个不同SM的单个延迟样本。皮尔逊相关性为0表示无相关性，1表示完全正（线性）相关，-1表示完全负相关。

相关性应用示例：图3(a,b)中的延迟分布具有0.998的皮尔逊相关性，表明分布几乎相同。然而，图3(a, c)之间的相关性约为-0.365，表明相关性极小。

V100相关性热力图：V100的皮尔逊相关性热力图如图6(a)所示。每个轴显示不同的SM，根据其GPC ID分组。热力图显示GPC0中的SM与GPC0和GPC1中其他所有SM之间存在高相关性。总的来说，同一GPC内的SM表现出（近乎）完美的相关性，并且相邻的两个GPC组具有非常高的正相关值。然而，距离较远的GPC相关性较低。该分析表明，两组GPC（GPC0&1和GPC4&5）具有相似的特性，与另外两个GPC（GPC2&3）相比。这很可能是由于GPU物理布局的对称性，GPC2&3位于中央，而GPC0&1和GPC4&5位于边缘（图4）。


图 6: (a) V100, (b) A100, 和 (c) H100 GPU上L2延迟的皮尔逊相关性热力图。对于A100和H100，GPC 0-3位于左侧“分区”，而GPC 4-7位于右侧“分区”。

观察 #4: 利用皮尔逊相关性的片上延迟分布相似性可以用来确定SM的布局。

C. GPU 分区的影响

A100的延迟特性变化：A100的皮尔逊相关性热力图（图6(b)）显示，单个GPC内的SM之间存在一些差异，这与V100中GPC内所有SM之间皮尔逊相关性相对恒定的情况不同。由于GPU的尺寸从V100增加到A100，无论是在物理芯片尺寸还是每个GPC的SM数量上，GPU实际上由两个“分区”组成，这影响了延迟特性。虽然未显示，但在V100上观察到的非均匀L2延迟在A100和H100中也存在。然而，A100的GPC之间的SM不一定具有相同的延迟分布特性（如在V100中观察到的那样）。热力图中的对角线显示，同一GPC内的SM之间的延迟几乎相同，但在V100中存在的相邻GPC的相似性在A100中减弱了。

H100的新层级CPC：如图6(c)所示，H100中GPC内SM之间的差异更大。一组4或6个SM（即2或3个TPC）具有相似的延迟特性，这与其他GPC内的SM组非常不同。这种延迟分布的相似性表明，在TPC和GPC层级之间可能存在一个新的核心层级，我们称之为CPC。

H100的SM间通信网络：GPC内CPC层级的高级框图如图7(a)所示，一个GPC内有三个CPC，它们通过一个SM-to-SM网络互连。与前几代GPU相比，H100的一个独特特性是分布式共享内存【12】，它能够通过SM-to-SM网络访问远程共享内存。为了解SM-to-SM网络的延迟特性，我们使用了一个类似于算法1的合成内核，但关键区别在于加载操作是针对同一GPC内的远程共享内存发出的。

SM-to-SM延迟特性：图7(b)显示了GPC内不同CPC组合之间SM-to-SM通信的平均延迟，以（源，目的）CPC配对信息表示。CPC的相对距离影响延迟——例如，当两个SM在CPC0内通信时延迟最低（196个周期），而对于可能距离最远的CPC2中的SM之间的通信，延迟增加到约213个周期。其他源和目的CPC的组合显示出基于其距离和位置的不同延迟。


图 7: (a) H100 GPC内提供SM间通信的CPC层级框图，以及(b) 基于CPC位置的平均SM间延迟。

观察 #5: 更大的GPU（具有多个“分区”）会导致不同的延迟特性和更强的非均匀性。特别地，H100在TPC和GPC之间增加了一个新的SM核心层级，这影响了延迟特性。

跨GPU分区的L2延迟：与V100相比，近期GPU（A100, H100）在跨越GPU“分区”时，L2延迟特性有显著不同。在图8(a,b,c)中，我们绘制了3种不同GPU的平均L2访问延迟，该延迟是GPC内所有SM到单个MP内不同L2 slice请求的平均值。V100的平均延迟约为212个周期（图8(a)），与本节前面的结果相似。然而，对于A100，平均L2延迟差异显著——例如，来自GPC 1,2,6,7的请求延迟与V100相似，但来自GPC 0,3,4,5的L2访问延迟要高得多。在我们的评估中，我们假设目的地（或L2 slice）位于左侧“分区”。由于GPC 1,2,6,7与目的L2 slice位于同一分区，当访问近分区时延迟较低。然而，对于GPC 0,3,4,5，目的L2 slice位于远分区——导致延迟高得多，约为400个周期，这是由于跨越GPU分区来回传输的额外延迟造成的。

H100的延迟特性：H100也像A100一样有两个分区。然而，其延迟值在各GPC间更为均匀（图8(c)）。虽然具体细节尚不清楚，但H100的内存层级（和缓存组织）进行了优化，其中“L2...缓存来自与该分区直接相连的GPC中SM的内存访问数据”【12】。因此，来自SM的延迟无论内存地址如何都相对相似。

L2未命中惩罚分析：为了更好地理解L2延迟行为，我们测量了L2未命中惩罚——即缓存行从主内存加载到L2所需的时间。对于V100和A100，L2未命中惩罚相对恒定（图8(d,e)），但对于H100，延迟不是恒定的，因为未命中惩罚会根据数据缓存的位置而变化（图8(f)）。本工作并未完全逆向工程GPU NoC架构或内存层级；然而，随着GPU尺寸的增加，引入了更多的延迟非均匀性。


图 8: 从一个GPC到一个MP的平均L2命中延迟（上图）和从全局内存到L2的平均L2未命中惩罚或延迟（下图），分别在(a,d) V100, (b,e) A100, 和 (c,f) H100上。

观察 #6: 带有多个GPU“分区”的近期GPU根据L2缓存策略会产生额外的延迟非均匀性。

IV. 片上网络带宽分析

在本节中，我们分析GPU的片上带宽，特别是到L2 slice（或内存分区）的带宽，即L2 fabric带宽。

A. 均匀的片上带宽

聚合L2带宽与内存带宽比较：在图9(a)中，我们比较了不同GPU上测量的总聚合片上L2 fabric带宽和全局内存带宽。聚合L2 fabric带宽是指所有SM向所有L2 slice发送流量（并在L2中命中），而内存带宽测量则假设访问未在L2中命中。两种带宽都是通过顺序内存访问测量的，内存带宽达到了峰值最大内存带宽的约85-90%。虽然片外内存带宽持续增加，但片上L2带宽也相应增加，比片外内存带宽高约2.4倍至3.5倍。尽管拥有更高的L2带宽并不令人意外【22】，但这确实对GPU NoC的设计有影响，我们将在第六节讨论，因为NoC显著影响L2 fabric带宽。


图 9: (a) 片上聚合L2和全局内存带宽。(b) V100中当单个SM访问时，到单个L2 slice的L2带宽分布。(c) V100中当单个GPC访问时，到单个L2 slice的L2带宽分布。

观察 #7: 在现代GPU中，片上聚合（总）L2 fabric带宽超过了总片外内存带宽。

带宽的均匀性：虽然片上延迟是非均匀的（观察#3），但到不同目的地（或L2 slice）的带宽则更为均匀。图9(b)显示了V100中当一个SM向单个L2 slice发送流量时的L2带宽直方图，展示了所有源（SM）和目的地（L2 slice）组合的带宽分布。从单个SM到L2 slice的平均带宽约为34 GB/s，标准差σ = 0.147 GB/s。当一个GPC内的所有SM向单个L2 slice发送流量时，可实现的带宽约为85 GB/s，标准差σ = 0.06 GB/s，如图9(c)所示——这是在所有GPC和目的L2 slice组合中测量的。这两个结果都显示了带宽的紧密分布，并表明从不同核心（SM）和GPC提供给每个L2 slice的带宽大致相似。结果还表明，要饱和单个L2 slice的带宽，至少需要4个SM。

观察 #8: 虽然到不同L2 slice的延迟是非均匀的，但到不同L2 slice的带宽（大部分）是均匀的。

输入加速比（Input Speedup）：在本文中，输入加速比定义为提供到GPU互连中的额外带宽【5】。GPU是分层组织的，输入加速比对于最小化分层带宽共享可能带来的性能瓶颈是必要的。在我们评估的GPU中，两个SM被集群到一个TPC中，TPC加速比（图11）实际上是TPC输出带宽与SM输出带宽的比率。在2对1共享的情况下，TPC加速比为1意味着两个SM总共只能提供50%的带宽。TPC加速比值为2则意味着两个SM可以向目的L2 slice的NoC提供“全”带宽。加速比是通过x个SM与一个SM的带宽（性能）比率来测量的，其中SM向所有L2 slice发送流量。对于TPC加速比，x=2；对于GPCl，x是TPC的数量；对于GPCg，x是TPC数量乘以每个TPC的SM数量。


图 11: 说明GPU NoC内提供的各种互连加速比的框图。

不同GPU的加速比表现：图10显示了不同GPU在读和写操作上的加速比。对于读操作，从SM所需的带宽（即小的请求消息）低于到SM的带宽（即大的带有缓存行数据的回复消息），而写操作则相反。对于TPC加速比，所有GPU的读操作都提供了全带宽（即TPC加速比为2）；然而，对于写操作，V100上仅实现了1.09倍的加速比。对于A100和H100，TPC写操作也实现了全带宽。对于GPCl，要实现全带宽，V100、A100和H100分别需要7、8和9的加速比。V100达到了这个加速比的约50%，而H100接近85%，加速比接近8。评估表明，提供了输入带宽加速比以最小化分层组织对性能的影响，并且近期的GPU拥有更多的片上带宽或输入加速比。特别是，GPCg相比GPCl提供了额外的加速比，尤其是在A100和H100上。对于H100，也显示了CPC加速比，结果表明，增加的层级对读访问没有影响，但当CPC内的所有SM都用于写操作时，由于只提供了约4.6的加速比（而全带宽需要6的加速比），性能略有下降。


图 10: 不同GPU架构的互连输入加速比比较。

观察 #9: 尽管GPU是分层组织的，但存在输入加速比以向片上互连提供更高的带宽。

B. GPU分区带来的非均匀带宽

跨分区访问的带宽非均匀性：在近期的具有多个GPU分区的GPU（即A100, H100）中，跨分区时引入了额外的非均匀延迟（观察#6），但在本节中，我们展示了在访问不同分区中的L2 slice时如何出现带宽非均匀性。A100的L2带宽非均匀性的一个例子如图12所示。SM0与L2 slice 0-39位于同一分区，而SM2与L2 slice 40-79位于同一分区。对于“近”分区，可以实现更高的带宽（约39.5 GB/s），但对于到“远”分区的L2访问，带宽下降到约26 GB/s。对于位于另一分区的SM2，可以实现类似的L2带宽，但“近”和“远”L2 slice的位置是相反的。


图 12: A100上(a) SM0和(b) SM2的L2 slice带宽，其中SM0和SM2位于不同的GPU分区。

L2带宽分布：图13(a)显示了A100中当每个SM访问单个L2 slice时的带宽分布。对于其他L2 slice也显示了类似双峰分布的结果——较高的峰值出现在与L2 slice在同一分区的SM，较小的峰值出现在远程分区的SM。图13(b)显示了H100的L2带宽分布，由于本地化的L2缓存【12】，它呈现出单峰。与V100相比，A100和H100的每个L2 slice带宽更高，但变化也更大。


图 13: (a) A100和(b) H100 GPU中到单个L2 slice的L2带宽分布。

远程带宽较低的原因：为了理解在A100中访问远程分区L2时带宽较低的原因，图14显示了随着更多SM向同一L2 slice发送流量，L2 slice带宽如何增加。在大约8个SM时，无论L2是在近分区还是远分区，带宽都达到饱和。然而，对于较少数量的SM，远分区可实现的带宽较低——即，对于一或两个SM，远分区可实现的带宽最多低28%。由于到远分区L2 slice的延迟较高，根据利特尔定律（Littel’s Law），较长的延迟导致较低的带宽，直到使用足够的SM来饱和带宽。


图 14: A100上随着发送流量的SM数量增加，单个L2 slice的带宽变化。

观察 #10: 近期GPU（例如A100, H100）相比之前的GPU拥有更多的片上带宽（和更多的加速比）；然而，这也导致了非均匀的带宽。

L2输入加速比：由于单个内存分区（MP）内有多个L2 slice，进入MP的带宽可能成为瓶颈，需要来自NoC输出的L2输入加速比（图11）。为了评估L2输入加速比，我们改变了L2 slice的数量（及其在MP间的分配），并评估一个所有SM都在发送流量的内核。Contiguous MP表示使用的L2 slice在同一个MP内，而Distributed MP表示L2 slice分散在不同的MP中——因此，对于Contiguous MP，访问的MP数量不变，但对于Distributed MP，MP数量从1增加到4。

L2输入加速比的评估：结果显示，在V100上，两种方法之间的带宽（性能）差异极小（图15(a)）——这表明存在近乎理想的L2输入加速比，以最小化L2 slice共享NoC输出带宽带来的影响。


图 15: V100上的L2带宽，(a) 所有SM访问不同数量的L2 slice，(b) 不同数量的SM访问一个MP，以及(c) 14个SM访问不同数量的MP。

GPC加速比的影响：相比之下，在向一个MP发送流量时改变SM的位置会导致显著的性能下降（图15(b)）。Contiguous SM将所有SM放在1个GPC（对于14个SM）或2个GPC（对于28个SM）中，与将SM分布在6个GPC的Distributed SM相比，性能显著下降——当28个SM访问单个MP时，下降约62%——这是由有限的GPC加速比引起的。为了进一步理解GPC加速比的影响，使用了14个SM，它们的流量被分配到不同数量的MP。对于Distributed SM，对性能的影响相对较小；然而，对于Contiguous SM，当MP目的地数量从1增加到4时，性能增加了218%——因此，这表明GPC加速比的一部分是在空间上提供的（即，额外的连接性），而不仅仅是在时间上（即，更多的带宽）。

观察 #11: 由于NoC加速比的不对称性，将内核负载均衡到各个SM比负载均衡到L2 slice更为关键。

C. 片上流量

片上流量与地址哈希：GPU的片上流量由内存流量（或L2流量）决定，特别是内存地址。如果一种内存访问模式导致某个内存通道被过度使用，整体内存吞吐量（和系统性能）可能会下降——这种行为被称为内存露营（memory camping）【41】。为了防止内存露营，现代GPU采用复杂的地址哈希机制，以确保对所有L2 slice的访问是负载均衡的【42】。

真实工作负载的流量模式：为了理解哈希的效果，图16显示了来自Rodinia基准套件【43】的两个工作负载（bfs, gaussian）在V100 GPU上的内存访问模式，该图是时间的函数，不同的强度表示流向每个L2 slice的流量大小。尽管内存访问量随时间变化，但流量在各个通道间的分布保持相对一致。因此，从NoC角度对片上流量进行负载均衡不是关键，但需要为每个目的地提供均匀的带宽以支持地址哈希。


图 16: 在NVIDIA V100 GPU上测量的(a) BFS和(b) Gaussian工作负载的内存访问模式。

观察 #12: 由于内存流量在内存通道间是负载均衡的，因此NoC流量也是负载均衡的。

A7 补充细节

V. 对GPU时序侧信道攻击的影响

本节探讨NoC特性如何影响GPU时序侧信道攻击，以及非均匀延迟如何被利用于发起时序侧信道攻击，并可作为防御时序侧信道攻击的机制。

A. GPU协同定位

利用NoC特性进行协同定位：在高度并行的GPU架构中，侧信道（或隐蔽信道）攻击通常需要将内核精确地协同定位在核心（或SM）上，以利用GPU的并行性。然而，鉴于现代GPU中侧信道攻击的威胁日益增加，GPU供应商正在限制一些被用于侧信道攻击的性能计数器的可用性——例如，为回应先前的工作【44】，性能计数器现在没有root权限已无法使用【45】。此外，从分析器获得的信息也受到限制——例如，在近期的NVIDIA GPU中，每个L2 slice的性能计数器已不再可用【37】，只有聚合的性能计数器值可用。在本文中，我们论证了NoC延迟特性如何被利用来确定核心（SM）的布局信息，并可用于内核的协同定位。如前述观察#3所述，L2缓存访问延迟可作为估算SM与L2 slice之间距离的可靠代理。因此，可以利用皮尔逊相关性（观察#4）来逆向工程GPU内的核心布局。这包括确定SM在GPC内的位置以及L2 slice在内存分区内的位置。因此，在需要协同定位进行GPU侧信道攻击时【46】–【48】，可以利用这些延迟特性。此外，SM布局可以在GPU NoC输入端建立一个隐蔽信道，但如果希望在GPU NoC的输出端（或L2的输入端）建立隐蔽信道，L2 slice的布局也可能被利用。

启示 #1: NoC特性分析可以被利用来获取关于核心和/或L2 slice布局的信息，用于GPU的时序侧信道攻击。

B. 攻击：非均匀延迟

非均匀延迟对时序攻击的影响：性能或执行时间与输入相关的微架构特性之间的关联，是GPU时序侧信道攻击中普遍利用的手段【6】，【7】。时序受L2延迟影响，对于给定的源（SM）和目的地（L2），延迟不会改变，但如果内核在不同核心上执行，延迟或时序就会改变。在本节中，我们探讨NoC延迟对两种先前已提出的GPU时序侧信道攻击——AES密钥恢复【6】和RSA侧信道攻击【49】——的影响，并讨论NoC对攻击和防御机制的启示。

1) AES密钥恢复：先前的工作【6】利用了唯一缓存行请求数量与执行时间之间的线性关系来恢复GPU中的AES密钥。对于使用一个warp（32个线程）的合成内核，根据内存合并的程度，延迟与访问的唯一缓存行数量成线性比例（图17(a)）。这些结果与先前的工作【6】，【7】一致，后者已在不同GPU上展示了类似的线性关系。然而，这种关系仅在利用相同的SM（或核心）时成立。当分配不同的SM而目的地（或L2 slice）不变时，线性关系会“平移”，如图17(a)中不同SM所示。由于威胁模型假设攻击者知道这种线性关系，攻击者会获取使用可能密钥的AES内核的执行时间。然而，除非攻击者确切知道内核的核心位置，否则唯一缓存行的范围可能会有很大差异。例如，如果观察到240个周期的延迟，唯一缓存行的数量可能从12到18不等。因此，当AES内核在不同核心上执行时，必须知道特定核心的线性关系；否则，攻击者无法使用时序准确推断出唯一缓存行请求的数量并猜测密钥字节。


图 17: (a) 当唯一缓存行数量变化时的平均计时（延迟）和(b) 在A100上不同核心（SM）的square内核平均执行时间。

2) RSA密钥恢复：先前的工作【49】展示了如何根据RSA密钥中1的数量与其执行时间的线性关系来推断其比例。在解密过程中，RSA对位0执行square()和reduction()，对位1执行额外的multiply()和reduction()。由于每个函数需要恒定的计算量，处理位1的执行时间大约是处理位0的两倍。因此，攻击者可以利用RSA程序执行时间随执行函数数量线性增加的特点来推断1的数量。然而，根据执行代码的核心（SM），由于到内存分区的非均匀延迟，内核执行时间可能会有很大差异。图17(b)显示了A100上不同SM上square()内核计时的一个例子。在这个例子中，SM 2, 16, 30, 44, 58位于一个分区，其他SM位于另一个分区。square内核【50】利用两个SM，为了便于说明，一个SM是固定的，而另一个SM的位置是变化的。根据其他SM的布局，计时（延迟）可能会有很大差异——即，当SM位于不同的GPU分区时，分区之间的额外延迟和同步开销导致执行时间更长，最多可达1.7倍。即使对于位于同一分区内的SM，非均匀延迟也引入了高达12%的性能差异。

新的攻击可能性：非均匀延迟是否可以被利用于不同类型的攻击还有待观察。例如，近期的工作【51】利用了多跳（2D mesh）网络中的距离和更高的延迟来确定L2访问模式。GPU中的跳数不一定不同，但延迟特性可能被利用来发动新型的侧信道攻击。

启示 #2: 非均匀L2延迟不会从根本上改变时序侧信道攻击，但核心之间的时序差异会影响侧信道攻击的行为。


图 18: GPU上的AES密钥恢复攻击，(a) 传统（静态）和(b) 随机线程块调度。图中圈出了正确的密钥值，为简化起见，只显示了16个密钥中的4个。

C. 防御：随机线程块调度

随机线程块调度：虽然先前的侧信道攻击已在真实GPU系统上得到验证，但它们需要收集大量数据样本，而且每次执行工作负载时，核心（或SM）的分配都是相同的——即，由于分配了相同的SM，观察到的延迟是相同的——因此没有观察到非均匀延迟。GPU中的线程块（或CTA）调度负责将内核分配给核心，实际上是静态调度。因此，我们提出使用随机分配的动态调度，即在每次内核执行时，线程块最初被映射到随机的SM，以防止侧信道攻击。随机线程块调度实际上是一种随机种子调度，即核心的分配从一个随机的核心开始。

1) AES密钥恢复：随机线程块调度对AES密钥恢复的影响如图18所示，图中x轴为不同的可能密钥值，y轴为皮尔逊相关性值。皮尔逊相关性用于比较时序和唯一缓存行请求数量的分布，当猜对密钥字节时，会获得一个高相关性值。使用传统的静态调度，图18(a)显示了如何恢复AES密钥。然而，使用我们提出的随机线程块调度，高皮尔逊相关性消失了（图18(b)），因为随机调度导致AES的执行时间基于到L2缓存slice的非均匀延迟而变得不确定。


图 19: RSA时序攻击，(a) 静态和(b) 随机线程块调度。

2) RSA密钥恢复：随机线程块调度对RSA密钥恢复的影响如图19所示。使用静态线程块调度时，‘1’的数量与执行时间之间存在线性关系；然而，当RSA代码在不同SM上执行时，线性关系会根据核心位置而平移。相比之下，当使用随机线程块调度时，线性关系变得非常嘈杂，如图19(b)所示。RSA内部的单个内核计时根据布局而变化，导致整体执行时间存在显著的可变性。因此，如果攻击者有一个测量的计时，‘1’的数量范围可能会有很大差异——例如，执行时间为$2 \times 10^9$个周期时，‘1’的数量可能在416到1920之间变化。

启示 #3: 随机线程块调度可以利用非均匀延迟来减少GPU中时序侧信道攻击的影响。

随机化作为防御手段：随机化常用于防止在侧信道攻击中收集有用的时序测量数据【52】–【54】。然而，随机化通常被添加到硬件中（例如，随机化合并逻辑【52】），这会引入一些复杂性并可能降低性能。相比之下，我们的随机线程块调度实际上是一种随机种子调度，不会引入额外的硬件复杂性。虽然随机线程块调度可以减少先前提出的时序侧信道攻击的影响，但该方案是否能完全防御各种时序侧信道攻击还有待观察。

VI. 对GPU NoC架构的影响

在本节中，我们探讨NoC分析对GPU NoC本身的影响。更重要的是，我们重新审视了先前GPU NoC研究【28】中确定的GPU NoC瓶颈，并展示了先前假设的NoC瓶颈如何未能正确模拟NoC带宽对整体内存和L2带宽的影响。

A. NoC中的带宽瓶颈

NoC流量模式与瓶颈识别：GPU中的片上网络（NoC）由一个请求网络和一个回复网络组成（图20(a)），其流量模式为“多对少对多”【28】——其中少数的内存控制器相对于众多的核心成为性能瓶颈。由于高吞吐量工作负载的内存访问通常是读操作，请求大小较小但回复数据较大，因此回复带宽（图20(b)中的BWnoc−mem）被认为是吞吐量处理器的NoC瓶颈【28】，【30】，【31】，【33】。当回复网络的NoC-MEM接口处出现拥塞时，拥塞背压不仅影响内存带宽利用率，还影响请求网络。


图 20: 吞吐量处理器中“多对少对多”的通信模式。(a) 突出了对分带宽（$BW_{NoC−Bc}$），(b) 突出了内存带宽（$BW_{MEM}$）和接口带宽（$BW_{NoC−MEM}$）。

仿真与真实系统的差异：然而，我们重新审视了回复带宽瓶颈，并评估了其对内存带宽利用率的影响。使用与先前工作【28】相似的配置，图21(a)绘制了在执行一个内存密集型合成内核时，单个内存通道的内存利用率随时间的变化。内存带宽达到了最大带宽，但该带宽无法持续，平均内存带宽利用率显著降低——约为20%——因为回复带宽接口拥塞阻塞了内存系统并向请求NoC施加背压。然而，对真实GPU的评估，如前文图9(a)所示，在不同GPU上都提供了高带宽利用率，超过了85%。因此，模拟器的NoC建模不一定能代表真实的GPU系统。


图 21: 基于仿真的评估中，由回复带宽瓶颈的背压导致的内存通道利用率波动。

NoC对L2带宽的影响：除了内存带宽，L2 fabric带宽或提供给L2的片上带宽也受到NoC带宽的影响。如前文图9所示，近期GPU的L2带宽超过了整体内存带宽；然而，使用相同配置的模拟也显示L2带宽低于最大内存带宽。因此，GPU中的NoC需要被设计（和建模）成既不瓶颈内存带宽，也不瓶颈L2带宽。

启示 #4: 在真实GPU中，NoC不会成为内存（或L2）带宽的瓶颈，但不准确的NoC假设会限制整体系统带宽。

B. 带宽层级

带宽层级设计原则：先前工作【28】中吞吐量处理器的基线NoC在内存带宽和NoC对分带宽之间提供了平衡，以实现成本效益（或“吞吐量效益”）的NoC。然而，对分带宽（图20(a)中的$BW_{NoC−B}$）仅在节点（即请求网络的计算节点和回复网络的内存节点）注入足够带宽以饱和对分带宽时才成为一个重要指标。由于通信发生在核心和内存节点之间（反之亦然），来自节点的终端或接口带宽（特别是图20(b)中的$BW_{NoC−MEM}$）成为瓶颈。接口带宽不足从根本上限制了内存（和L2）带宽。

瓶颈分析：另一种看待这个问题的方法是使用瓶颈分析【55】来识别接口带宽可能是瓶颈。由于核心、NoC和内存系统是串联连接的，“K个串联子系统的最大吞吐量是子系统吞吐量的最小值”【56】。因此，即使提供了高内存带宽，如果聚合终端带宽（$BW_{NoC−MEM}$）成为瓶颈，系统也无法维持如此高的带宽，即使提供了足够的对分带宽。总的来说，片上带宽层级【57】应反映层级化的系统组织，内存带宽是成本最高但最有限的资源，而NoC带宽是（相对）成本最低的资源，应被配置为不瓶颈内存或L2带宽。

启示 #5: 为了确保NoC不成为整体系统的瓶颈，需要考虑片上带宽层级；特别是，除了对分带宽外，接口带宽也需要被恰当配置。

C. GPU中的基线NoC

基线NoC在仿真中的重要性：已经有许多基于仿真的GPU架构研究，包括NoC架构、warp调度和内存层级。对于这类评估，需要恰当地定义基线NoC，以确保得出正确的见解。在图22中，我们收集了来自几项先前工作【14】，【15】，【17】，【28】–【32】，【58】，【59】的数据。我们绘制了内存带宽（BW）与所使用的NoC-MEM接口带宽的关系，其中接口带宽定义为$BW_{NoC−MEM} = f_{NoC} \times w \times C$，（$f_{NoC}$: NoC时钟频率，w: NoC通道宽度，C: MP数量）。对于斜线下方的点（即$BW_{NoC−MEM} < BW_{MEM}$），NoC接口带宽是瓶颈，对于内存密集型工作负载可能限制整体性能，并实际上造成一个“网络墙”。虽然改善NoC带宽的不同技术可以提高这类系统的整体性能，但基线NoC或片上带宽成为系统瓶颈，而不是内存带宽。相比之下，当$BW_{NoC−MEM} > BW_{MEM}$时，内存带宽是瓶颈，这更能代表现代GPU系统。因此，GPU架构的模拟应确保片上带宽不会瓶颈更昂贵的全局内存带宽或造成“网络墙”——否则，得出的见解可能导致夸大NoC优化的好处。


图 22: 不同符号代表不同基于仿真的先前工作【14】, 【15】, 【17】, 【28】– 【32】, 【58】, 【59】中建模的内存带宽和NoC-Mem接口带宽的比较。

多跳拓扑的挑战：像GPU这样的吞吐量处理器对带宽敏感，应为不同目的地或内存分区提供均匀的带宽（第四节-A）。然而，如果假设多跳网络如2D mesh作为NoC拓扑，提供均匀带宽是一个挑战，如图23(a)所示。根据核心节点的物理位置，一些节点将获得高得多的吞吐量，最多可达2.4倍。如果使用全局公平仲裁，如基于年龄的仲裁【61】，【62】，那么可以在2D mesh拓扑上实现全局公平带宽（图23(b)），但代价是流控制的额外复杂性。此外，在像2D mesh这样的扁平拓扑中提供任何必要的加速是一个挑战。


图 23: 多跳（2D mesh）拓扑架构中出现的带宽不公平现象，(a) 轮询仲裁和(b) 全局公平的基于年龄的仲裁。

现代NoC拓扑假设：近期的工作【34】，【35】已利用层级交叉开关作为基线NoC拓扑，它不一定有2D mesh NoC那样的限制。虽然层级交叉开关可以提供更多带宽，但在进行GPU架构探索时，需要仔细配置NoC到内存（和L2）的带宽。最近一个GPU模拟器探索（Accel-Sim【36】）讨论了L2带宽差异如何可能解释其在建模现代GPU时的一些不准确性。因此，在模拟中需要考虑和恰当建模整个系统的带宽层级，而不仅仅是NoC带宽。

启示 #6: 多跳拓扑（例如，2D mesh）在为所有节点提供均匀带宽（无论其位置如何）以及在NoC中提供所需加速方面都存在挑战。

A4 实验环境

A4 实验结果

1. 片上延迟特性 (第III节)

2. 片上带宽特性 (第IV节)
* 均匀性: 与延迟相反，从SM到不同L2 slice的带宽在V100上是基本均匀的，带宽分布非常集中 (图9b, 9c)。
* 带宽层级: 片上L2 fabric的总带宽远高于片外内存带宽（约2.4x-3.5x），表明NoC为L2提供了充足的带宽 (图9a)。GPU通过输入加速比（Input Speedup）来避免层级结构带来的瓶颈 (图10)。
* GPU分区影响: A100在跨分区访问时，带宽也呈现非均匀性，访问远分区的带宽显著低于近分区（约39.5 GB/s vs 26 GB/s）。这是因为长延迟限制了少量SM的带宽，需要约8个SM才能饱和远端带宽 (图12, 图14)。H100由于L2缓存的本地化，带宽分布更均匀 (图13)。
* 负载均衡: 由于GPC到NoC的加速比有限，将计算任务（SMs）均匀分散到不同GPC比将内存访问（L2 slice）分散更重要 (图15)。
* 流量模式: 现代GPU的地址哈希机制能有效将真实应用的内存流量均匀分布到所有L2 slice，因此NoC的流量也是负载均衡的 (图16)。

3. 对时序侧信道攻击的影响 (第V节)
* 攻击面: NoC的非均匀延迟特性可被攻击者利用，通过测量延迟来推断内核所在的SM物理位置，辅助攻击 (启示#1)。
* 攻击干扰: 这种非均匀性也干扰了现有的时序攻击（如AES, RSA密钥恢复），因为执行时间会因内核被调度到不同SM而变化，使得固定的“时间-操作”关系不再可靠 (图17, 启示#2)。
* 防御机制: 利用非均匀延迟，可以通过随机线程块调度来防御时序攻击。每次执行时将内核随机分配到不同SM，引入时序噪声，破坏攻击者建立精确时序模型的能力 (图18, 图19, 启示#3)。

4. 对NoC架构设计的影响 (第VI节)
* 瓶颈再评估: 实验证明，真实GPU的NoC带宽充足，并不会成为内存或L2带宽的瓶颈，这与一些早期基于仿真的研究结论（即“回复网络瓶颈”）相反 (图21 vs 图9a, 启示#4)。
* 设计原则: GPU NoC设计应遵循带宽层级原则，即NoC接口带宽必须被充分配置以支持更昂贵的内存和L2带宽，而不仅仅是关注对分带宽 (图20, 启示#5)。
* 仿真假设: 许多仿真研究中对NoC带宽的假设过低，可能导致错误的结论（如“网络墙”），并高估NoC优化的效果 (图22)。2D-mesh等多跳拓扑因难以提供均匀带宽和加速比，不适合作为现代GPU NoC的基线模型 (图23, 启示#6)。

A5 结论

本文对现代GPU的片上网络（NoC）在延迟和带宽方面进行了详尽的分析。研究发现，基于核心和L2 slice的物理位置，存在非均匀的延迟，而片上带宽则近似均匀。然而，具有多个“分区”的大型GPU在带宽和性能方面引入了额外的非均匀性。基于对NoC的特性分析，我们介绍了一些其带来的影响，包括对GPU时序侧信道攻击以及NoC架构设计的影响。