【CUDA调优指南】缓存&访存流程

前言

学习 UP 主 比飞鸟贵重的多_HKL 的 【CUDA调优指南】缓存&访存流程 视频，记录下个人学习笔记，仅供自己参考😄

refer 1：【CUDA调优指南】缓存&访存流程

refer 2：https://chatgpt.com/

1. L1/L2 Cache

这节我们来讲 ncu 与 kernel 调优指南第二讲—缓存&访存流程，主要是来了解 GPU 中的各级缓存以及访存的主要流程，后面会简单介绍 Memory Chart 和 Memory Chart 中的各个指标及其含义，主要参考的是 NVIDIA 官方文档：https://docs.nvidia.com/nsight-compute/NsightCompute

这节我们并不会有相关代码调优分析，更多的是对上节合并访存的一个补充

我们先看下 NVIDIA 官方文档对于 Caches 的描述：

数据从全局内存（global memory）到 SM 的传输过程中，会先去 cache 中查询是否有缓存，如果 cache 命中，则下一次的访问时间会大幅减少。在启动了 L1 缓存的情况下，对全局内存的访问将经过 L1 Cache；如果未命中，则会接着从 L2 缓存中查找；如果再次未命中，则会从全局内存 DRAM 中读取

上图给出了 NVIDIA GPU 中 L1TEX 级别的缓存其内部各单元如何协同完成请求分发、命中/未命中判断、与 L2 二级缓存交互以及最终把数据送回给 SM（Streaming Multiprocessor）的流水线模型

下面我们分段来讲下其中各条数据/请求的流向：(from ChatGPT)

1. SM 侧如何发起请求

• Warp Scheduler → MIO
  • SM 里的多个 warp 调度器把它们的 load/store（全局 global、局部 local 和共享 shared）或者纹理/Surface 请求发送到一个叫 MIO（Memory I/O crossbar）的模块
• MIO → L1TEX
  • 全局（global）和局部（local）内存访问，连同 shared（片上共享内存）访问，一起通过 MIO → LSUIN
  • 纹理和 Surface 内存访问则通过 MIO → TEXIN

2. L1TEX 内部关键阶段

2.1 Tag Stage（标签阶段）

• 只针对 global/local 的 load/store 以及所有 texture/surface 请求做 Tag 查找
• 命中（hit）：把对应 cache line（sector）送到 Data Stage
• 未命中（miss）：把缺失的 sector 请求送到 Miss Stage 去向下游（L2）取数

2.2 Miss Stage（未命中阶段）

• 汇总所有未命中扇区（sector），把它们打包后发给 L2
• 当 L2 返回数据时，再把扇区交给 Data Stage

2.3 Data Stage（数据阶段）

• 接收来自 Tag Stage（命中）或 Miss Stage（回填）的扇形数据
• 根据请求类型，分发到 LSU Data（普通 Load/Store 路径）或 TEX Data（纹理采样数据）

2.4 LSU Data

• 把数据送回给 SM 的 Warp Scheduler，用于普通的 GPU load/store

2.5 TEX Data + TEX Filter

• 先把扇区原始字节送入 TEX Data，再通过 TEX Filter 做插值/过滤（linear，point，shadow sampling 等）
• 最终返回给 SM，用作纹理采样的颜色或深度值

2.6 Shared 内存请求

• 直接从 LSUIN 绕过 Tag/Miss 阶段，走专门的数据通路到 LSU Data，因为 shared memory 并不走 L1 cache 的常规标签匹配

3. 与 L2 Cache 的交互

• Miss Stage → L2
  • 所有 L1 未命中的扇区统一发给 L2
• L2 → L1（Miss Stage/Data Stage）
  • L2 将扇区数据返回给 Miss Stage，再流入 Data Stage 分发给 LSU 或 TEX

总结下整体流程：

• 双输入：Load/Store 请求（LSUIN）和 Texture 请求（TEXIN）同时竞争 L1TEX 的标签和数据通路
• 三级流水：Tag Stage 判定 → Miss Stage（若未命中）→ Data Stage 汇聚
• 双输出：LSU Data（给普通 Load/Store）和 TEX Data+Filter（给纹理采样）两条回路
• Shared 内存：特殊绕过路径，效率更高
• 多级缓存协作：L1TEX 负责快速命中，Miss 时下探到 L2，再回填 L1，以提高整体带宽与吞吐

我们再来看看 L2 Cache：

上图展示了 GPU 中二级缓存（L2）的逻辑结构和数据/控制流，说明所有上层缓存（如 L1TEX、L1V etc.）以及其他内存客户端如何通过 L2 与主存（DRAM）交互，以及 L2 内部如何处理命中、未命中、原子操作和一致性请求

值得注意的是 NVIDIA GPU 上的 L2 并不是一个单一、巨大的存储体，而是被分成若干个“切片”（Slice），每个切片负责一部分地址范围，这些切片物理上分布在芯片各处，访问时可以并行处理

在官方文档中，整个 L2 Cache 称作 ltc（L2 cache），而单个切片就用 lts（L2 cache slice）来区分。每个 lts 切片内部又包含 Tag Stage（lts_t）、Miss Stage（lts_m）和 Data Stage（lts_d）等子模块

图中画的那个绿色大框标注为 LTS 就是在讲一个 L2 Cache Slice 的内部流程，而不是把所有切片合起来的整个 L2。整个 L2（ltc）是 SM 共享的资源，但在硬件实现和性能度量上，它又被拆分成多个可并行的切片（lts），方便更细粒度地观察瓶颈

下面我们分段来讲下其中各条数据/请求的流向：(from ChatGPT)

1. 整体定位

• L2 Cache 位于芯片片上各缓存客户端和物理内存之间，统一处理所有来自各个 SM 的缓存未命中、DMA、纹理写回、帧缓冲读写等请求
• 物理地址空间：L2 以物理地址为索引，负责各 SM 之间以及 CPU/Host 与 GPU 之间的缓存一致性和数据共享
• 附加功能：除了基本的读写缓存，L2 还集成了
  • 压缩/解压缩 硬件（减少对外存带宽压力）
  • 全局原子操作单元（实现 GPU 级别的 atomicAdd、CAS 等原子操作）

2. 请求的入口与标签判断

2.1 来自 L1 的 Miss Stage

• 所有 L1（包括 L1D、L1C、L1TEX 等）未命中的扇区请求，会被汇总发给 L2 的 Tag Stage

2.2 Tag Stage（标签阶段）

• 在物理地址空间查找对应扇区是否已缓存在 L2
• 命中（hit）：标签匹配，产生 “hit” 扇区
• 未命中（miss）：标签未匹配，产生 “miss” 扇区，并将这些 miss 通知给 Data Stage 去下探主存

3. 数据阶段与主存交互

3.1 Data Stage（数据阶段）

• 接收来自 Tag Stage 的两类扇区
  • hit 扇区：直接从 L2 的 Data RAM 读出
  • miss 扇区：打包后下发给相应的内存控制器（Mem Controller），去 DRAM 读取
• 返回：DRAM 取回的数据先进 Data Stage，再回填到 L2 Data RAM，同时送回给请求的 L1

3.2 一致性（Coherence）

• L2 连接芯片内部的高速互联（L2 fabric），负责跨多 SM 的缓存一致性请求（如其它 SM 对同一扇区的写入、失效通知等）
• 一致性消息也是作为扇区级的请求进 Data Stage，由硬件自动处理

4. 全局原子操作

• 对于需要硬件支持的 全局原子指令（atomicAdd、atomicCAS 等），L1 扇区会被路由到 L2 的 Atomic 单元：
  • 原子指令操作数（operand）送入 Atomic 单元
  • 在 L2 Data RAM 上原地更新并返回结果
  • 保证跨 SM 的原子性与一致性

5. 压缩/解压缩（可选）

• 当内存带宽成为瓶颈时，L2 可对写入扇区做 硬件压缩，对读取扇区做 硬件解压
• 从而在 “miss 下探” 或 “回填” 阶段减少物理 DRAM 的数据量，提高带宽利用率

总结下整体流程：

• 统一入口：所有 L1 Miss、DMA、帧缓冲访问、纹理写回等都通过 L2
• 标签判断：Tag Stage 快速区分命中/未命中
• 数据流：命中走内部 RAM，未命中打包送主存，返回后回填与送返
• 一致性：借助 L2 fabric 实现多 SM 之间缓存一致性
• 原子操作：内置 Atomic 单元，直接在 L2 Data 上完成全局原子指令
• 压缩加速：可选的压缩/解压硬件，缓解外存带宽压力

OK，以上就是关于 L1/L2 Cache 访存的一个完整流程分析

关于更多细节的描述大家可以参考：https://docs.nvidia.com/nsight-compute/ProfilingGuide/index.html#hardware-model

2. Memory Chart

Memory Chart（内存图）以图形化逻辑方式显示 GPU 上和 GPU 外内存子单元的性能数据，包括传输大小、命中率、指令或请求数等

上篇文章我们简单分析了下合并访存代码的内存图，下面我们来跟随官方文档一起来看看 Memory Chart 的详细分析：

Note：以下内容均来自于 NVIDIA 官方文档：https://docs.nvidia.com/nsight-compute/ProfilingGuide/index.html#memory-chart

2.1 Logical Units (green)

逻辑单元在图中显示为绿色

• Kernel：在 GPU 流式多处理器（SM）上执行的 CUDA kernel
• Global：CUDA global memory
• Local：CUDA local memory
• Texture：CUDA texture memory
• Surface：CUDA surface memory
• Shared：CUDA shared memory
• Load Global Store Shared：指令直接将数据从全局内存加载到共享内存，无需经过中间寄存器文件访问

2.2 Physical Units (blue)

物理单元在图中显示为蓝色

• L1/TEX Cache：L1/Texture cache，底层物理内存由该缓存和用户管理的共享内存组成
• Shared Memory：CUDA 用户管理的 shared memory，底层物理内存由该内存和 L1/TEX Cache 组成
• L2 Cache：L2 cache
• L2 Compression：L2 Cache 的内存压缩单元
• System Memory：片外 system (CPU) memory
• Device Memory：执行 Kerenl 的 CUDA 设备的片上 device (GPU) memory
• Peer Memory：其他 CUDA 设备的片上 device (GPU) memory

Note：根据 GPU 架构的不同，图中所示单元的具体组合可能会有所不同，因为并非所有 GPU 都拥有所有单元

2.3 Links

Kernel 与其他逻辑单元之间的连线表示针对相应单元执行的指令（Instruction）数量，例如，Kernel 和 Global 之间的连线表示从全局内存空间加载数据或存储数据到全局内存空间的指令。如果是 NVIDIA A100 的 GPU 则 Load Global Store Shared 模式的指令将单独显示，因为它们的寄存器或高速缓存访问行为可能不同于常规的全局加载或共享存储

逻辑单元（绿色）和物理单元（蓝色）之间的连线表示因各自指令而发出的请求（Req）数，例如，从 L1/TEX Cache 到 Global 的连线显示了全局内存加载指令产生的请求数

每个连线的颜色代表相应通信路径峰值利用率的百分比，图右侧的颜色图例展示了从未使用（0%）到峰值性能（100%）的渐变。如果某条链路不活跃，则以灰色显示，图例左侧的三角形标记对应图中的各条链路，相较于单纯的颜色渐变，这些标记可以更准确地估算所达到的峰值性能

2.4 Ports

一个单元通常会共用一个数据端口，用于输入和输出流量。尽管共享该端口的各条链路可能远未达到它们各自的峰值性能，但该单元的数据端口可能已达到峰值利用率。端口利用率在图中以位于输入和输出链路处、单元内部的彩色矩形表示。端口使用与数据链路相同的颜色渐变，并在图例左侧配有相应的标记，不活跃的端口则显示为灰色

2.5 Metrics

可以使用 --set full、--section MemoryWorkloadAnalysis_Chart 或 --metrics group:memory__chart 命令行收集该图表的各项指标。下图展示了内存表（Memory Tables）中报告的峰值与内存图表（Memory Chart）中端口之间的对应：

3. Memory Tables

上面我们简单分析了内存图（Memory Chart），下面我们来看看内存表（Memory Tables）

内存表展示了各种内存硬件单元的详细指标，例如共享内存、各级缓存和设备内存。对于大多数表格条目，你可以将鼠标悬停在单元格上以查看底层指标名称和描述，有些条目是从其他单元格计算得出的派生指标，本身不显示指标名称，而是显示相应的计算结果。如果某个指标未参与通用派生计算，则在工具提示中显示未 UNUSED，你还可以将鼠标悬停在行或列标题上，查看表格这一部分的说明

Note：以下内容均来自于 NVIDIA 官方文档：https://docs.nvidia.com/nsight-compute/ProfilingGuide/index.html#memory-tables

3.1 Shared Memory

Columns：

• Instructions：对于每种访问类型，每个 warp 实际执行的所有汇编（SASS）指令总数
• Requests：共享内存的所有请求总数。在 SM 7.0（Volta）和更新的架构上，每条共享内存指令都会产生一个请求
• Wavefronts：为响应所请求的共享内存数据所需的 wavefronts 数量，wavefronts 会被串行化，并在不同的时钟周期中处理
• % Peak：峰值利用率的百分比。数值越高意味着该单元的利用率越高，可能反映潜在的瓶颈，但并不一定表示使用效率也高
• Bank Conflicts：如果多个线程请求的地址映射到同一内存 bank 的不同偏移量，这些访问就会被串行化。硬件会根据需要将冲突的内存请求拆分成多个相互无冲突的独立请求，从而减少有效带宽，减少的系数等于相互冲突的内存请求数

这里博主对 Wavefronts 这个指标有些困惑，它代表什么含义呢？为什么表格中的值刚好是 Instructions/Requests 的 32 倍，是巧合还是说有什么联系呢？🤔

Wavefronts 在 Shared Memory 表里，其实就是 处理一整组线程（即一个 warp）所花费的 SMEM 周期数。通俗的讲：(from ChatGPT)

1. 什么是一个 Wavefront？

在 NVIDIA GPU 上，一个 warp 包含 32 个线程，编译器会把它当作最小的执行单元。NCU 里称它为 wavefront，意思就是 “同时被调度、一起访问内存的那 32 个线程”

2. Requests vs. Wavefronts

• Requests：表示你的 kernel 有多少次 “warp 级别” 的 SMEM 访问指令发出
• Wavefronts：表示实际上 SMEM 总共启动了多少个 warp 来处理这些访问
• 如果访问完全没有 bank 冲突，每个 warp 级访问只要一个周期就能搞定，那么 Wavefronts == Requests

3. 为什么表里 Wavefronts = Requests x 32？

其实不是巧合，而是 kenrel 中那段访问模式出现了 最严重的 bank 冲突。SMEM 有 32 个 bank，如果同一个 warp 里 32 个线程都去读同一条 bank，就只能 串行 处理，每次只能搞定 1 条线程的请求，要分 32 个周期才能把整个 warp 的请求都处理完

• 冲突时，每多一个线程跟它抢 bank，就多一个额外的 wavefront 周期，NCU 把多出来的周期记到 Bank Conflicts 列
• 公式上，Wavefronts = Requests + BankConflicts
• NCU 测到的数值正好是：Wavefronts = 32 768 + 1 015 808 = 1 048 576 = 32 768 × 32

简单总结下：

• Wavefronts = 实际启动的 warp 数（含因 bank 冲突而串行的额外次序）
• Requests = 理论上 warp 级的访问次数（不算冲突）
• 当极端冲突时，Wavefronts = Requests × warp_size（32），这正是表里看到的 32 倍关系，绝非纯粹巧合，而是因为每个 warp 的 32 条线程都抢同一 bank 导致了完全串行的访问

Rows：

• (Access Types)：共享内存访问操作（load/store/atomic 等）
• Total：同一列中所有访问类型的总和。

3.2 L1/TEX Cache

Columns：

• Instructions：对于每种访问类型，每个 warp 实际执行的所有汇编（SASS）指令总数
• Requests：每种指令类型生成的所有 L1 请求总数。在 SM 7.0（Volta）和更新的架构上，每条指令都会为 LSU traffic（global、local…）生成一个请求，而对于纹理（TEX）traffic，则可能生成多个请求。在示例中，65536 条全局加载指令中的每一条指令都会产生一个请求
• Wavefronts：为响应所请求的内存操作所需的 wavefronts 数量，wavefronts 会被串行化，并在不同的时钟周期中处理
• Wavefronts % Peak：处理 wavefronts 的单元峰值利用率百分比。数值过高可能意味着处理流水线已饱和并可能成为性能瓶颈
• Sectors：发送到 L1 的所有 L1 扇区访问总数。每个加载或存储请求都会访问 L1 缓存中的一个或多个扇区。原子和归约操作会传递到 L2 缓存
• Sectors/Req：L1 缓存的 sectors 与请求的平均比率。在相同数量活跃线程的 warp 中，比率越小意味着内存访问模式越高效。对于拥有 32 个活跃线程的 warp，每个访问大小的最优比率如下：32-bit: 4, 64-bit: 8, 128-bit: 16。较小的比率表明在同一缓存行内存在一定程度的访问均匀性或加载重叠；较高的比率则可能意味着非合并（uncoalesced）的内存访问，从而导致内存流量增加
  • 在示例中，全局加载的平均比率为每次请求 32 个 sectors，这意味着每个线程需要访问不同的 sector。理想情况下，对于拥有 32 个活跃线程的 warp，如果每个线程访问单个对齐的 32 位值，那么比率应为 4，因为每 8 个连续的线程会访问同一个 sector
• Hit Rate：L1 缓存中的 sector 命中率（未丢失的请求 sector 百分比）。发生未命中的 sector 需要向 L2 发出请求，从而增加对 L2 的 sector 未命中次数。更高的命中率意味着更好的性能，因为访问延迟更低，请求可以直接由 L1 响应，而无需进入后续阶段。不要将其与 Tag Hit Rate 混淆
• Bytes：从 L1 请求的字节总数。其大小等于 sector 数量乘以 32 byte，因为 L1 的最小访问大小为 1 个 sector
• Sector Misses to L2：在 L1 中未命中并在 L2 缓存中产生后续请求的 sector 总数。
  • 在示例中，全局和本地加载的 262144 次 sector 未命中，可通过将 12.5% 的未命中率乘以 2097152 个 sector 来计算得出
• % Peak to L2：用于发送二级缓存请求的 L1-to-XBAR 接口的峰值利用率百分比。如果该数值较高，则工作负载可能由写入、原子、归约这些操作主导，这会增加延迟并导致 warp stalls
• Returns to SM：从 L1 缓存发回 SM 的返回数据包数量。请求访问大小越大，返回的数据包数量越多
• % Peak to SM：XBAR-to-L1 返回路径的峰值利用率百分比。如果这个数字较高，则工作负载可能以读取操作为主，从而导致 warp stalls。通过提高读取操作合并或 L1 命中率可以降低这种利用率

Rows：

• (Access Types)：各种访问类型，例如从全局内存加载数据或在 surface 内存上的归约操作
• Loads：同一列中所有 load 访问类型的总和
• Stores：同一列中所有 store 访问类型的总和
• Total：同一列中所有 load 和 store 访问类型的总和

3.3 L2 Cache

Columns：

• Requests：对于每种访问类型向 L2 缓存提出的请求总数。该指标与 L1 缓存的 sector 未命中到 L2 相对应，每个请求最多可访问单个 128 字节高速缓存行中的四个扇区
• Sectors：对于每种访问类型从二级缓存请求的 sector 总数。每个请求访问 1~4 个 sector
• Sectors/Req：L2 缓存的 sectors 与请求的平均比率。在相同数量活跃线程的 warp 中，比率越小意味着内存访问模式越高效。较小的比率表明在同一缓存行内存在一定程度的访问均匀性或加载重叠；较高的比率则可能意味着非合并（uncoalesced）的内存访问，从而导致内存流量增加
• % Peak：占峰值持续扇区数的百分比。二级缓冲中的 “work package” 就是一个 sector。数值越高意味着单元利用率越高，可能反映潜在的瓶颈，但并不一定表示使用效率也高
• Hit Rate：L2 缓存中的 sector 命中率（未丢失的请求 sector 百分比）。发生未命中的 sector 需要向更后续的阶段发出请求，从而增加 Sector Misses to Device/System/Peer 中的一种。更高的命中率意味着更好的性能，因为访问延迟更低，请求可以直接由 L2 响应，而无需进入后续阶段。
• Bytes：从 L2 请求的字节总数。其大小等于 sector 数量乘以 32 byte，因为 L2 的最小访问大小为 1 个 sector
• Throughput：二级缓存吞吐量，单位为字节/秒。数值越高表示单元利用率越高
• Sector Misses to Device：L2 中未命中但在设备内存中产生后续请求的 sector 总数
• Sector Misses to System：L2 中未命中但在系统内存中产生后续请求的 sector 总数
• Sector Misses to Peer：L2 中未命中但在 peer 内存中产生后续请求的 sector 总数

Rows：

• (Access Types)：各种访问类型，例如来自 L1 cache 的加载或归约操作
• L1/TEX Total：来自 L1 缓存的所有操作总数
• ECC Total：由 ECC（Error Correction Code）引起的所有操作的总和。如果 ECC 使能，部分修改 sector 的 L2 写请求会导致从 DRAM 加载相应的 sector。这些额外的加载操作会增加 L2 的 sector 未命中次数
• L2 Fabric Total：连接两个 L2 分区的 L2 Fabric 上的所有操作总数。仅显示在带 L2 Fabric 的 CUDA 设备上 kernel 执行时的情况
• GPU Total：二级缓存所有客户端的所有操作总数

3.4 L2 Cache Eviction Policies

Columns：

• First：使用 evict_first 策略在二级缓存中访问的 sector 数。使用该策略缓存的数据将在 eviction 优先级顺序中排在第一位，并有可能在需要 eviction 缓存时被驱逐。该策略适用于流式数据
• Hit Rate：使用 evict_first 策略访问二级缓存中 sector 的缓存命中率
• Last：使用 evict_last 策略在二级缓存中访问的 sector 数。使用该策略缓存的数据在 eviction 优先级顺序中将排在最后，可能只有在其他使用 evict_normal 或 evict_first 策略的数据被驱逐后才会被驱逐。该策略适用于需要在缓存中持久保留的数据
• Hit Rate：使用 evict_last 策略访问二级缓存中 sector 的缓存命中率
• Normal：使用 evict_normal 策略在二级缓存中访问的 sector 数。这是默认策略
• Hit Rate：使用 evict_normal 策略访问二级缓存中 sector 的缓存命中率
• Normal Demote：使用 evict_normal_demote 策略在二级缓存中访问的 sector 数
• Hit Rate：使用 evict_normal_demote 策略访问二级缓存中 sector 的缓存命中率

Rows：

• (Access Types)：各种访问类型，例如来自 L1 cache 的加载或归约操作
• L1/TEX Total：来自 L1 缓存的所有操作总数
• L2 Fabric Total：连接两个 L2 分区的 L2 Fabric 上的所有操作总数。仅显示在带 L2 Fabric 的 CUDA 设备上 kernel 执行时的情况
• GPU Total：二级缓存所有客户端的所有操作总数

3.5 Device Memory

Columns：

• Sectors：对于每种访问类型，从设备内存请求的 sector 总数
• % Peak：设备内存使用峰值百分比。数值越高意味着该单元的利用率越高，可能反映潜在的瓶颈，但并不一定表示使用效率也高
• Bytes：在二级缓存和设备内存之间传输的字节总数
• Throughout：设备内存吞吐量，单位为字节/秒。数值越高表示该单元利用率越高

Rows：

• (Access Types)：设备内存加载和存储操作
• Total：同一列中所以访问类型的总和

4. add1 核函数缓存 & 访存分析

前面我们跟随 NVIDIA 官方文档简单学习了 L1/L2 Cache 以及 NCU 分析后提供的 Memory Chart 和 Memory Tables

下面我们就简单看下我们自己写的 add1 核函数经过 NCU 分析后的内存图表，代码如下：

void __global__ add1(float* x, float* y, float* z){
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    z[idx]  = x[idx] + y[idx];
}

Memory Chart 如下：

下面我们按图把 add1 这个最简单的加法 kernel 的访存流程分几步再来走一遍：

1. 全局内存指令数 vs. 请求数

• 全局内存指令（Global Inst） 总数：786.43 K
  • 其中 524.29 K 是 load 指令
  • 262.14 K 是 store 指令

2. L1/TEX Cache 完全失效

• 图中 L1/TEX Cache 的 Hit Rate: 0.00%
• 意味着全局的 load/store 操作 没有一笔 命中 L1，全部都得 “穿透” 到下一级，也就是 L2 Cache

3. L1 → L2：Misses 转发量

• 只看 load 请求（524.29 K 次），每次访存按 128 B 四个 sector 发出：$524.29\text{ K req}\times 128\text{B}\approx 67.11\text{ MB}$

4. L2 Cache：33.37% 命中

• L2 在这 67.11 MB 的 sector 请求里，命中了 33.37 %，剩余未命中的请求需要下钻到 DRAM 中

5. L2 → Device Memory：Misses 处理

• 剩下在 L2 也 miss 的请求被转发到片外 DRAM（Device Memory）

接着来看下 Memory Tables，先看下 L1 Cache：

这个 L1/TEX Cache 表其实把 “L1 级别” 的访存行为拆得非常细，我们可以从中读出以下关键信息：

1. 各类指令 & 请求数量

• Global Load
  • Instructions（warp 级）：524 288
  • Requests（coalesced warp 访问数）：524 288
• Global Store
  • Instructions：262 144
  • Requests：262 144
• 其它种类（Local、Texture、Surface、Shared 等）都是 0，说明这个核函数只做了纯粹的全局 load/store

2. Wavefronts = Requests

• Load 对应 524 288 个 wavefront（每个 warp 对 x/y 各发一次 load）
• Store 对应 262 144 个 wavefront
• 且二者都一一对应（Wavefronts==Requests），说明在 L1/TEX 层 没有 bank 冲突，每个 warp 访问都能一次性并发发出并占用一个周期

3. Coalescing 情况：Sectors/Req = 4

• 表中 Sectors 列：
  • Load 拿到了 2 097 152 个 sector
  • Store 写入了 1 048 576 个 sector
• 而 Sectors/Req = 4，说明每个 warp 级请求实际跨了 4 个 sector
• 从 Bytes 列可以看出：
  • Load：2 097 152 x 32 B/sector = 67.11 MB
  • Store：1 048 576 x 32 B/sector = 33.55 MB
• 这里的 sector 大小是 32 B，所有每个 warp 访问了 4x32 B = 128 B，正好跟一个 warp（32 条 4B float）完全对齐—典型的 “完美合并”

4. 完全没命中 L1 → 全部 Miss 转给 L2

• Hit Rate = 0 %：Load/Store 在 L1/TEX 全部 Miss
• Sector Misses to L2 = Sectors 数量（Load+Store 共 3 145 728），也就是说所有的 sector 请求都要到 L2

5. 对 L1/TEX 硬件带宽的利用率

• % Peak（峰值利用率）
  • Load 用了 4.63%
  • Store 用了 2.32%
  • 合计 6.95%
• 说明这个 kernel 里 L1/TEX 并没有被跑满

6. 后续到 L2 的流量 & 返回

• Percent Peak to L2：Load 占 L2 峰值带宽的 18.53%，Store 9.26%，合计 ~27.8%
• Returns to SM（回到 SM 的 sector 数）共 786 432

简单总结下：

• 这个最简单的 add1 核，完全没用到 L1/TEX 缓存（0% hit），所有全局读/写都到 L2
• Warp 级 coalescing 很理想（4 sectors/req = 128 B），而且无 bank 冲突（Wavefronts == Requests）
• L1/TEX 层利用率很低（~7%），L2 又只能在命中三分之一，剩下都跑到 DRAM

通过这张表，我们可以非常细粒度地看到每次 warp 访问在 L1 发生了什么、拿了多少 sector、冲不冲突、用了多少带宽，以及 miss 后打到 L2 的流量情况

最后我们来看下 L2 Cache 和 Device Memory：

从上面这张 L2 Cache + Device Memory 的表里，我们可以看出 add1 在二级缓存和片外内存上的整个数据流特征：

一、L2 Cache 层级

1. Load 全部 Miss → 524 288 次 warp-load 共 2 097 152 扇区（≈67 MB）都打到了 DRAM

2. Store 全部 Hit → 262 144 次 warp-store 共 1 048 576 扇区（≈33 MB）都在 L2 就被接过，不算 Miss，但后续会以“正常驱逐”形式刷到 DRAM

3. 整体命中率 33.33 % 正好等于 store 在所有访问（524 288+262 144=786 432）中的比例

4. L2 带宽利用率 ~45.6 %：30.4 % 用于 load，15.2 % 用于 store

另外，L2 的 Eviction Policy 一栏也能看到：

• Load miss 全部走 “Normal” 驱逐（2 097 152 扇区）。
• Store hit 也走 “Normal” 驱逐（1 048 576 扇区），但命中率 100 %。

二、Device Memory 层级

1. DRAM 读流量 ≈67 MB（2 097 160 扇区），写流量 ≈32 MB（1 013 704 扇区），合计 ≈99.5 MB。

2. 带宽利用率高达 92.8 %，说明整个 kernel 已几乎把片外内存带宽跑满了。

3. 读带宽约 216 GB/s，写带宽约 105 GB/s，总计 ~321 GB/s。

简单总结下：

• add1 的 load 全部越过 L2（0% 命中），直接从 DRAM 取回，成为主导流量
• store 虽然在 L2 命中，但最终仍以驱逐方式写入 DRAM，也贡献了 ~32 MB 的写流量
• 二级缓存自身只用了 ~45 % 带宽，而片外内存几乎用满 (~93 %)，表明此 kernel 最主要的瓶颈在 DRAM 带宽。如果能提高 L2 对 load 的命中，比如通过数据重用或分块，就能显著减轻 DRAM 压力、提升性能

OK，以上就是通过 NCU 分析 add1 这个简单 kernel 访存流程的全部内容了

Note：关于其它 kernel 我们上篇文章也简单分析过，这边博主就不再赘述了

结语

这篇文章我们主要是了解了 GPU 中的各级缓存以及访存的主要流程，并通过查阅 NVIDIA 官方文档学习了 Memory Chart 和 Memory Tables 中的各项指标及其含义

最后，我们从一个简单的合并访存核函数出发，利用 NCU 分析了其访存流程

OK，以上就是本篇文章的全部内容了，大家可以多看看 NVIDIA 的官方文档，还是非常有收获的🤗

下载链接

• CUDA 调优指南代码下载链接【提取码：1234】

参考

• 【CUDA调优指南】缓存&访存流程
• https://docs.nvidia.com/nsight-compute/NsightCompute