【CUDA调优指南】合并访存

前言

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

refer 1：【CUDA调优指南】合并访存

refer 2：https://chatgpt.com/

1. 合并访存的概念

在 CUDA 中，合并访存（memory coalescing） 是指同一个 warp（通常为 32 个线程）内的线程对全局内存（global memory）进行访问时，如果它们访问的地址是连续且对齐的，硬件就可以把多次分散的小访问合并成一次或少数几次大的内存事务（transaction），以此来最大化带宽利用率、减少延迟，并显著提升性能

Note：内存事务（memory transaction）是指硬件层面在访存总线上一次性装载或写回的一段连续字节

什么是合并访存？

• Warp 访问模式：CUDA 以 warp（32 线程）为单位调度和访存
• 连续 & 对齐：如果线程 t 访问地址 A + t * sizeof(元素类型)，并且起始地址 A 对应内存事务大小（通常是 128 字节）对齐，那么硬件只需发起一次内存事务，就能将整个 warp 的请求打包完成
• 硬件优化：合并访存减少了访存事务的数量，从而提高全局内存带宽利用率，降低访存延迟，提升整体吞吐

合并访问的优点包括：

• 带宽利用率更高：少发事务、每次传输更多数据
• 减少访存延迟：从几十甚至上百个小事务，变成少数几次大事务
• 更高的吞吐率：当访存不再成为瓶颈，GPU 的计算单元可以被更充分地利用

下面我们举个简单的例子来说明下：

假设我们有一个一维浮点数组 float* data，长度足够长，每个 block 有 128 个线程（一个 warp + 若干线程）

合并访存（Coalesced）示例代码如下：

__global__ void kernel_coalesced(const float* data, float* out) {// 全局线程索引int idx = blockIdx.x * blockDim.x + threadIdx.x;// 每个线程访问 data[idx] —— 完全连续out[idx] = data[idx] * 2.0f;
}// 典型调用
// blockDim.x = 128, gridDim.x = N
kernel_coalesced<<<gridDim, blockDim>>>(d_data, d_out);

• 访问模式：线程 0 访问 data[0]，线程 1 访问 data[1]，…，线程 31 访问 data[31]
• 事务合并：这 32 次 4 字节的访问可以被合并成一次 128 字节的读，再打包写回

非合并访存（Non-Coalesced）示例代码如下：

__global__ void kernel_strided(const float *data, float *out, int stride) {int idx = blockIdx.x * blockDim.x + threadIdx.x;// 每个线程访问 data[idx * stride] —— 间隔较大out[idx] = data[idx * stride] * 2.0f;
}// 假设 stride = 16
kernel_strided<<<gridDim, blockDim>>>(d_data, d_out, 16);

• 访问模式：线程 0 访问 data[0]，线程 1 访问 data[16]，线程 2 访问 data[32]，…
• 事务拆分：因为每次访问相隔 64 个字节（16x4 字节），一个 128 字节的事务只能覆盖 data[0…31]，data[16…47]，data[32…63] 等多个子区间，不同线程的请求就无法合并到同一次事务里，需要多次访存，严重浪费带宽

因此在实际 CUDA 优化中，尽量让数据布局和访问模式对齐并连续，这样才能充分发挥 GPU 全局内存带宽，减少访存瓶颈

2. 合并访存的代码实现

代码实现如下：

#include <bits/stdc++.h>
#include <iostream>
#include <cuda_runtime.h>
#include <random>
#include <ctime>
#include <sys/time.h>
#include <cublas_v2.h>void __global__ add1(float* x, float* y, float* z){int idx = blockIdx.x * blockDim.x + threadIdx.x;z[idx]  = x[idx] + y[idx];
}int main(){const int N = 32 * 1024 * 1024;float* input_x_host = (float*)malloc(N * sizeof(float));float* input_y_host = (float*)malloc(N * sizeof(float));float* input_x_device;float* input_y_device;cudaMalloc((void**)&input_x_device, N * sizeof(float));cudaMalloc((void**)&input_y_device, N * sizeof(float));cudaMemcpy(input_x_device, input_x_host, sizeof(float), cudaMemcpyHostToDevice);cudaMemcpy(input_y_device, input_y_host, sizeof(float), cudaMemcpyHostToDevice);float* output_host = (float*)malloc(N * sizeof(float));float* output_device;cudaMalloc((void**)&output_device, N * sizeof(float));dim3 grid(N / 256);dim3 block(64);for(int i = 0; i < 2; ++i){add1<<<grid, block>>>(input_x_device, input_y_device, output_device);cudaDeviceSynchronize();}cudaMemcpy(output_host, output_device, N * sizeof(float), cudaMemcpyDeviceToHost);free(input_x_host);free(input_y_host);free(output_host);cudaFree(input_x_device);cudaFree(input_y_device);cudaFree(output_device);return 0;
}

整个项目的目录结构如下：

cuda_learn/
├── CMakeLists.txt
└── kernel_profiling_guide├── CMakeLists.txt├── combined_access.cu└── ...2 directories, 3 files

cuda_learn 文件夹下的 CMakeLists.txt 内容如下：

cmake_minimum_required(VERSION 3.20.0)
project(cuda_practice VERSION 0.1.0 LANGUAGES CUDA CXX C)
find_package(CUDAToolkit)
add_subdirectory(kernel_profiling_guide)

kernel_profiling_guide 文件夹下的 CMakeLists.txt 内容如下：

add_executable(combined_access combined_access.cu)
target_link_libraries(combined_access PRIVATE CUDA::cudart ${CUDA_cublas_LIBRARY})
if(CMAKE_BUILD_TYPE STREQUAL "Debug")
target_compile_options(combined_access PRIVATE $<$<COMPILE_LANGUAGE:CUDA>:-G>)
endif()
target_compile_options(combined_access PRIVATE -lineinfo)/* add new */
// add_executable(transpose transpose.cu)
// target_link_libraries(transpose PRIVATE CUDA::cudart ${CUDA_cublas_LIBRARY})
// if(CMAKE_BUILD_TYPE STREQUAL "Debug")
//     target_compile_options(transpose PRIVATE $<$<COMPILE_LANGUAGE:CUDA>:-G>)
// endif()
// target_compile_options(transpose PRIVATE -lineinfo)

编译运行指令如下：

cd cuda_learn
mkdir build && cd build
cmake .. && make -j24
./kernel_profiling_guide/combined_access

OK，整个代码实现后我们就来分析下代码中具体做了些什么

3. 合并访存代码分析

合并访存的示例代码非常简单，它就是去计算 input_x 和 input_y 两个向量的和，最终把结果存储在 output 中

值得注意的是我们开启的线程总数只覆盖了 1/4 的数据：

dim3 grid(N / 256);
dim3 block(64);

虽然我们给 input_x、input_y 以及 output 分配的大小是 N * sizeof(float)，但是我们只处理了 8 * 1024 * 1024 大小的浮点数，并没有处理全部的数据，因此开启的线程总数只有原始 N 个 float 数据的 1/4

核函数实现如下：

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

代码也非常简单，首先获取这个线程的全局索引 idx，接着计算 x 和 y 的和并写入到 z 中，那 z[idx] = x[idx] + y[idx]; 这行代码包括两次全局内存的读取，一次全局内存的写入，它其实涉及到的东西挺多的

我们绘制了一个图来说明整个数据流向：

1. 参数与整体框架

数组长度 N = 8 * 1024 * 1024，约 32 MB 数据

图中绿色框表示逻辑单元（Logical Unit），它包括 Kernel 层面的指令和全局内存逻辑；蓝色框表示物理单元（Physical Unit），是指硬件上的 L1/L2 Cache 和设备 DRAM

2. Kernel→Global（逻辑层面）

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

• 每个元素做 2 次读 + 1 次写 ⇒ 总共 3N 次全局内存操作
• CUDA 以 warp（32 线程为一组）为最小调度粒度
  • warp 数量 = N / 32 = 8 * 1024 * 1024 / 32 = 262144 个 warp
  • 每个 warp 发出 3 条 “global load/store” 指令 ⇒ $\text{Inst}=3\times \frac{N}{32}=786432$

3. Global→L1（逻辑→物理请求转换）

Global 层面每条 load/store 指令，会在全局内存接口处转换成一次“请求”（Request）：

• Load 请求
  • 2 次 load/warp $\times$ 262144 warp = 524288 load Req
• Store 请求
  • 1 次 store/warp $\times$ 262144 warp = 262144 store Req

4. L1→L2（物理层面的 sector 事务）

在硬件层面，Global 接口的每次请求不是一次 4 Bytes 或 128 Bytes 传输，而是拆成若干个 32 Bytes 的最小事务单元（sector）：

• 每个 sector = 32 Bytes
• 因为我们的 32 个线程连续读 32 x 4 Bytes = 128 Bytes，且地址恰好 128 Bytes 对齐，所以每条 load 或 store 请求都被切分成 4 个连续的 32 Bytes sector

因此：

• L1→L2（load） 传输了 $524288\text{ Req}\times 4\text{ sector}/\text{Req}=2097152\text{ sector}$
• L1→L2（store） 同理 $262144\text{ Req}\times 4\text{ sector}/\text{Req}=1048576\text{ sector}$

Note：NVIDIA GPU 在全局内存访问时会以最小 32 字节为单位（sector）进行传输

5. L2↔Device Memory（DRAM 事务）

同样的每个 32 Bytes sector，若 L2 Cache 未命中，就沿着路径下钻到 DRAM（Device Memory），再把它拉回到 L2（有 L2 Cache） 和 L1（若启用了 L1 缓存），此时 L2↔DRAM 也会产生与 L1↔L2 同量级的 sector 事务

那大家可能会有所困惑，为什么我们的数据要经过 L1 Cache 和 L2 Cache 呢？为什么不直接访问 DRAM 呢？

目前我们的访问流程如下：

• 1. 线程发起全局内存读请求
• 2. L1 Cache
  • 如果开启了 L1 缓存，硬件会先在每个 SM 自带的 L1 cache 中查找
• 3. L2 Cache
  • 如果 L1 miss，就到片上共享的 L2 cache 查找，所有对全局内存的访问都必须先经过 L2
• 4. Device Memory（DRAM）
  • 如果 L2 也 miss，才发起对设备内存（global memory）的 DRAM 访问，完成后会将对应的 cache line 带回 L2 和 L1

全局内存就是上图中的设备内存（DRAM），而 L1/L2 则分别位于 SM（Streaming Multiprocessor）内部和整个芯片上的硬件缓存，可以用来加速对全局内存的访问

GPU 在物理上也遵循 “层级化内存”（memory hierarchy）的设计—从最快的小容量缓存，到最慢的大容量 DRAM。把数据先经过 L1/L2 Cache 能够将热点数据留在低延迟、高带宽的存储层，大幅减少对高延迟 DRAM 的请求并隐藏剩余延迟，从而保持高吞吐

add1 kernel 的访存模式是最理想的，线性、对齐、全 warp 连续访问，能够完美的合并访存

至此，我们的合并访存代码就分析完毕了，那你可能会有疑问，这么简单的程序要经过这么多的计算嘛，有没有什么好的工具能够帮助我们把数据流向以及数据大小等直接算出来，不用人工去计算呢，那这个时候就需要 Nsight Compute 了

4. Nsight Compute 分析合并访存核函数（add1）

NVIDIA Nsight Compute 是 NVIDIA 提供的性能分析工具，它专注于单个 CUDA 内核（kernel）的详细分析，可以提供指令级别的计数、内存吞吐量、线程执行、寄存器使用、共享内存利用率等详细指标

Nsight Compute 其实在安装完 CUDA 之后就已经安装了，具体安装位置在 /usr/local/cuda/bin，如下图所示：

Nsight Compute 提供两种界面：

1. ncu（命令行工具）：

• 用途：通过命令行运行 Nsight Compute，直接在终端中对 GPU 内核进行采样和分析
• 特点：适合自动化脚本、批量分析和生成报告文件（如 .ncu-rep），便于嵌入到开发和 CI/CD 流程中
• 使用方式：通过传递参数（例如 --launch-count）运行可执行文件，并将分析结果输出到终端或指定的报告文件中

2. ncu-ui（图形用户界面工具）

• 用途：提供直观的图形界面，用于加载和浏览由 ncu 生成的报告文件
• 特点：界面中会显示各种图表和 Breakdown 表格，如 GPU Throughput 图、Compute 和 Memory 的各项细分指标，便于交互地查看和分析性能数据
• 使用方式：启动 ncu-ui 后，通过菜单加载 .ncu-rep 报告文件，或直接配置目标程序进行交互式分析

关于 Nsight Compute 的简单使用大家可以参考：【模型分析】Nsight Compute使用入门

关于 Nsight Compute 更详细的介绍和使用推荐大家参考 NVIDIA 官方文档：https://docs.nvidia.com/nsight-compute/NsightCompute/#nsight-compute

下面我们就来利用 Nsight Compute 的 ncu-ui 图形化界面分析下 add1 核函数，看其数据流向以及数据大小是否和我们前面分析的结果一样

4.1 基础配置

首先终端执行 ncu-ui 启动 Nsight Compute 的图像化界面，如下图所示：

接着按照如下的步骤进行配置：

先点击上方菜单栏的 Connection 进行配置，接着点击 Linux (x86_64)，然后选择要分析的可执行文件的路径，这是关于上半部分 Target Platform 的配置

我们接下来看下半部分 Activity 的配置：

点击 Profile，选择 Common 下的输出文件（Output File）路径，Sections 下选择 full，最后点击右下角的 Launch 即可

执行过程中你可能会遇到如下的问题：

日志信息提示说我们没有权限访问 NVIDIA GPU 的性能计数器，此时我们可以通过 sudo 来访问，对于博主来说指令如下：

sudo /usr/local/cuda-11.6/bin/ncu-ui

接着再次进行上面的配置，点击 launch 后就会开始分析，分析完成后你会得到一份 .ncu-rep 的报告，如下所示：

4.2 Memory Chart

下面我们来具体分析下其中的一些细节，主要是来看 Memory Chart 内存图表的分析：

我们先点击 Page 中的 Details 详细信息页面，接着点击 Memory Workload Analysis 内存负载分析栏，选择 All，之后可以看到 kernel 执行过程中的内存图以及下面详细的表格

我们主要来分析其中的内存图：

Memory Chart 其实和我们前面自己绘制的那个图非常像，绿色代表逻辑单元，蓝色代表物理单元，我们简单对比下看我们之前分析的是否存在问题：

1. Kerenl → Global

• Memory Chart：786.43 K Inst
• 自绘图中：3·N/32 (warp) = 3×262 144 = 786 432 Inst

两者 786 K 基本吻合，就是每个 warp 发 2 个 load + 1 个 store 的请求

2. Global → L1

• Memory Chart：524.29 K Req（load）和 262.14 K Req（store）
• 自绘图中：524 288 load Req，262 144 store Req

两者完全一致

3. L1 ↔ L2：合并成 32 B-sector 事务

• 自绘图中：
  • Load：524 288 Req x 4 sector
  • Store：262 144 Req x 4 sector
  • sector 大小 32 B ⇒
    • 读：524 288 x 4 x 32 B = 67 108 864 B ≈ 67.11 MB
    • 写：262 144 x 4 x 32 B = 33 554 432 B ≈ 33.55 MB
• Memory Chart：
  • 从 L1 到 L2（以及相反方向）的箭头上标记：
    • 67.11 MB
    • 33.55 MB

这正是我们自己手算的 4-sectorx32 B 的读写总量

4. L1 命中率 = 0%

• Memory Chart：L1/TEX Cache Hit Rate: 0.00 %
• 含义：每次访问都是第一次读，L1 完全没有命中，所有请求都落到 L2

5. L2 ↔ Device Memory

• 自绘图中：假设 L2 无命中，所有 sector 都要下钻到 DRAM
  • 读：67.11 MB 下钻
  • 写：33.55 MB 写回
• Memory Chart：
  • L2 Cache Hit Rate = 33.37%
  • 从 Device Memory 到 L2 的箭头：67.11 MB（load miss）
  • 从 L2 到 Device Memory 的箭头：32.44 MB（store writeback）

注意这里写回 32.44MB 而不是 33.55 MB，说明一部分写操作可能被 L2 合并或优化掉了（L2 命中了部分 store 或做了 write combine），所以实际写会量略低于理论值

6. L2 命中意味着什么

• L2 命中率 ~33 %，说明大约三分之一的 load 请求在 L2 就满足了，不必再跑 DRAM；其余 ~67 % 仍然下钻，正好对应那 67.11 MB 的 load traffic

这里博主有个困惑，那就是既然 L2 命中率是 33% 则意味着 67.11MB 的读请求中约 1/3 的请求不用下钻到 DRAM 中，但是为什么 Memory Chart 图中 L2 和 DRAM 的 load 请求又是显示的 67.11MB 呢？

这里要把两个概念区分开来：(from ChatGPT)

1. “Transfer Size” 箭头

在 Nsight Compute 的 Memory Chart 里，箭头上的数字（比如 67.11 MB）就是简单地把 “请求次数 x 每次事务大小” 相乘，画到每一条通路上-不论这些请求最终是不是都要跑到 DRAM

• 我们看到的 67.11 MB，恰好是 $524288\text{ load requests}\times (4\text{sectors／req})\times (32\text{B／sector})=67.11\text{ MB}$
• 这个数字在 L1→L2 和 Device Memory→L2 两条箭头上都一模一样，因为它就是同一批“逻辑请求”（L1 miss 请求）量

2. L2 Hit Rate (33.37 %)

这个命中率告诉你，在那 524 288 次 4-sector 读请求里，有 33.37 % 在 L2 就被满足了，根本没必要再去 DRAM

• 所以物理上真正跑到 DRAM 的数据量应该是：$67.11\text{ MB}\times (1-0.3337)\approx 44.7\text{ MB}$
• 但 Chart 的 “Transfer Size” 模式并不会把箭头上的 67.11 MB 减去这 33.37%-它始终画出 “请求总量×事务大小”，而不去区分命中和未命中

因此，我们在 Chart 中看到的是逻辑请求而不是物理上真正跑到 DRAM 的数据量

7. 其它通道都是 0

• Local、Texture、Surface、Shared 等路径全是 0，表明这个 kernel 仅做 Global load/store，不涉及纹理或共享内存

这就是 Memory Chart 内存图表呈现的具体内容，通过该图我们大致能知道在 GPU 中各个数据的大小以及流向

4.3 Memory Tables

看完了 Memory Chart 我们再往下面翻一翻，看看 Memory Tables 具体的数值以及指标：

首先我们来看 Shared Memory：

add1 核函数没有利用共享内存进行任何的加载或存储操作，因此指令数全为 0，所有的读/写操作都发生在全局内存中（通过 L1/L2 Cache）

接着我们来看看核心的 L1 缓存的详细性能数据：

我们把表格中的数据单独拎出来：

Category	Values
Instructions	786,432 (total)
Requests	524,288 (load), 262,144 (store)
Wavefronts	524,288 (load), 262,144 (store)
% Peak	4.63% (load), 2.32% (store)
Sectors	2,097,152 (load), 1,048,576 (store)
Sectors/Req	4 (load), 4 (store)
Hit Rate	0%
Bytes	67,108,864 (load), 33,554,432 (store)
Sector Misses to L2	2,097,152 (load), 1,048,576 (store)
% Peak to L2	18.53% (load), 9.26% (store)
Returns to SM	786,432 (total)
% Peak to SM	6.95% (total)

分析：

1. 指令数和请求数

• 786,432 指令 和 786,432 请求，与 Memory Chart 中的请求数相符。
• 524,288 load 请求 和 262,144 store 请求与 Memory Chart 中的 524.29K 和 262.14K 请求数量一致。

2. 每请求的 sector 数

• 每个请求对应 4 个 sector，2,097,152 sectors / 524,288 requests = 4 sectors（读请求）和 1,048,576 sectors / 262,144 requests = 4 sectors（写请求）。这个也与 Memory Chart 中的值一致。

3. 命中率 0%

• 命中率为 0% 表明所有请求都没有在 L1 缓存中命中，意味着所有的请求都需要去 L2 缓存或 DRAM 获取数据。
• 在 Memory Chart 中没有显示 L1 缓存命中的数据，因为所有数据都需要进一步到 L2 去处理。

4. 请求的字节数

• 读取数据：524,288 * 4 * 32 B = 67,108,864 B ≈ 67.11 MB
• 写入数据：262,144 * 4 * 32 B = 33,554,432 B ≈ 33.55 MB
• 这些值与 Memory Chart 中的 67.11 MB（load）和 33.55 MB（store）对比吻合。

5. sector 到 L2 的未命中

• L1 到 L2 的 sector 未命中数：
  • 读请求：2,097,152 sectors
  • 写请求：1,048,576 sectors
• 这些值与 Memory Chart 中 L1/TEX Cache → L2 Cache 之间的 2,097,152 sector 和 1,048,576 sector 相符。

6. 返回 SM

• 返回 SM 的请求数：786,432（等于总的指令数），这个表明 L1 缓存并没有缓存任何数据，所有请求最终都返回到 SM

7. sector 的峰值与吞吐量

• 峰值（% Peak）较低：
  • Load：4.63%（说明 L1 缓存的利用率较低）
  • Store：2.32%（存储操作的利用率也较低）
• 吞吐量（Throughput）：
  • Load：216,536,912,751.68 B（这表明加载的吞吐量非常高）
  • Store：104,667,423,851.32 B（存储的吞吐量也较大）

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

L2 Cache 数据如下：

Category	Values
Requests	524,288 (load), 262,144 (store)
Sectors	2,097,152 (load), 1,048,576 (store)
Sectors/Req	4 (load), 4 (store)
% Peak	30.37% (load), 15.18% (store)
Hit Rate	33.37% (for both load and store)
Bytes	67,108,864 (load), 33,554,432 (store)
Throughput	216,536,912,751.68 (load), 108,268,043,366.03 (store)

分析：

1. 请求数

• 524,288 load 请求 和 262,144 store 请求 与之前 Memory Chart 中显示的相符（524.29 K 和 262.14 K）。

2. 每个请求的 sector 数

• 每个请求对应 4 个 32 B sector：2,097,152 sectors (load) / 524,288 req = 4 sectors。
• 同理，写操作也对应 4 个 sector：1,048,576 sectors (store) / 262,144 req = 4 sectors。

3. 命中率

• L2 的命中率 33.37% 表示 33.37% 的 load 和 store 请求直接命中 L2 cache，无需访问 DRAM。
• 在 Memory Chart 中，这一命中率与 L2 Cache Hit Rate（33.37%）匹配，表明大约三分之一的读写请求在 L2 中被直接满足

4. 字节量

• 总的传输量：67.11 MB 读请求和 33.55 MB 写请求，与 Memory Chart 中的 67.11 MB 和 33.55 MB 对应

L2 Cache Eviction Policies 数据如下：

Category	Values
First	0
Hit Rate	0%
Last	0
Normal	2,097,152 (for load), 1,048,576 (for store)
Normal Demote	0
Hit Rate	0%

分析：

• L2 Cache Eviction Policies 相关数据提供了缓存行的清除策略，分别是 First、Last、Normal 和 Normal Demote。
• Normal 对应正常访问的缓存行，所有的 load 和 store 请求都落在这里。具体的请求数：
  • 2,097,152 load 请求 和 1,048,576 store 请求。
• 命中率 0% 表示缓存没有显式的命中或数据迁移到 L2 以外的部分。

Device Memory 数据如下：

Category	Values
Sectors	2,097,152 (load), 1,048,576 (store)
% Peak	62.58% (load), 30.25% (store)
Bytes	67,109,120 (load), 32,438,528 (store)
Throughput	216,536,912,751.68 (load), 108,267,917,871.32 (store)

分析：

1. 请求数和 sector 数

• 由于 DRAM 是后备存储，L2 Cache miss 的请求会最终到达 DRAM，因此这里的传输量与 L2 Cache 传输量匹配。

2. 字节量

• 67.11 MB（load）和 32.44 MB（store）对应着 Memory Chart 中 L2 和 Device Memory 之间的传输数据。

3. 总吞吐量

• 加载和存储的数据量与 Memory Chart 的吞吐量一致，表明数据从 L2 缓存到 DRAM 的传输量非常大。

OK，以上就是 Memory Tabels 内存表的简单分析，我们从内存表中可以获取各个性能指标数据，有助于帮助我们分析 Kernel 的瓶颈

5. 非对齐但连续访问（add2）

前面我们提到的 add1 核函数是最理想的情况，带宽利用率最高，能实现完美的合并访存，下面我们来看另外几种情况

首先来看 add2 的实现：

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

• 与 add1 唯一不同的是：每个线程的访问下标都偏移了 +1
• 线程 t 访问的是 x[base + t + 1] 而不是 x[base + t]

这样会出现什么问题呢？在 add2 中，虽然每个线程访问的是连续的元素—线程 0 访问 x[1]，线程 1 访问 x[2]，…，线程 31 访问 x[32]，看上去似乎也是完美连续，但是它们并没有对齐

Note：对齐 vs. 连续

• 连续 意味着线程访问的地址是递增的
• 对齐 意味着这段连续地址段在内存中的起始位置正好落在硬件传输粒度（sector 或 cache line）的边界上

前面我们提到过在 CUDA 中，最小的传输粒度是 32 B（一个 sector），warp 里 32 个线程访问总共是 32 x 4B = 128 B 上的连续空间

• 完美合并访存（例如 add1）要求这 128 B 本身就 128 B 对齐，也就是它的起始地址 addr % 128 == 0
• 这样，硬件只需发起 4 次 32 B-sector 事务（4 x 32 B = 128 B）就能满足整个 warp

为什么 add2 会跨段？内存会不对齐呢？我们来简单分析下

在 add2 里，我们把索引 +1：

int idx = blockIdx.x * blockDim.x + threadIdx.x + 1;

假设某个 warp 的起始 blockIdx.x * blockDim.x 恰好是 0，那么这个 warp 实际要读写的元素是 [1 ... 32]， 对应的字节区间是：

[1×4 B, 32×4 B + 3] = [4 B, 131 B]

而 32 B-sector 的边界是：

[0–31], [32–63], [64–95], [96–127], [128–159], …

你会发现 [4-131] 这个区间 跨越了 5 个 sector：

• 1. [0-31]（实际用到 4-31）
• 2. [32-63]
• 3. [64-95]
• 4. [96-127]
• 5. [128-159]（实际用到 128）

因此，每个 load 或 store 都要拆成 5 个 sector 事务，而不是理想状态下的 4 个

这样会带来一些性能影响，包括：

• 事务数从 4 增到 5：每个 warp 的读写都要多一次 32 B 传输，带宽浪费 +25%
• 总传输率增加：原本 4 x 32 B =128 B，现在变成 5 x 32 B = 160 B
• 更高的延迟：多一次事务意味着更多的 DRAM 往返、更多的缓存压力

add2 虽然访问连续，但对齐被 +1 打乱了，于是完整的 128 B 区间跨了 5 个 32 B sector，导致非理想的合并访存

Nisght Compute 分析结果如下：

相比于 add1，add2 的 Memory Chart 有以下几点显著差异：

1. 扇区数（sectors）增加，缓存层间流量放大

在 add1 中，一个 warp 连续且对齐的 128 B 区域只跨 4 个 32 B sector；而 add2 因为在起始地址上 “+1” 字偏移，导致同样的 128 B 区间跨越了 5 个 sector，结果是从原本约 67 MB 的 L1→L2 传输量增加到约 75 MB。

2. L1 Cache 命中首次非零

由于多个 warp 访问的 5-sector 区间出现了部分重叠，add2 在 L1 层面产生了约 15.8% 的命中（add1 为 0%），这在一定程度上减少了下钻到 L2 的请求数，也让 L2↔SM 返回的数据量从 67 MB 降至约 40 MB。

3. L2 Cache 命中率提升

更高的 L1 命中加上扇区重用，使得 add2 的 L2 Cache 命中率从 33.4% 提升到 41.7%。不过，虽然命中率提高，但整体对齐和连续性的破坏仍然导致更大的跨层带宽开销，影响了总体的内存访问效率。

Memory Tables 的具体指标如下图所示，这边博主就不带着过了，大家可以自己看看：

此外 ncu-ui 还给了我们几条建议，我们来简单看下：

这几条建议核心都是围绕 “减少每次请求的扇区数／缓存行数”，具体要点如下：

1. L1TEX 层面的 Load/Store 合并

• 现在每个 Warp Load 平均用了 4.6 个 32 B 扇区（146.5 B），Store 用了 5.0 个扇区（160 B），理想情况都应是 4.0×32 B = 128 B。
• 要检查 Source Counters，修正线程间的 stride 或对齐方式，使全局读写在 L1TEX 只跨正好 4 个扇区。

2. L2 层面的 Load/Store 合并

• L1TEX 到 L2 每次请求是一个 128 B Cache line，但 Store 只用了 2.7/4 个扇区，Load 只用了 2.9/4 个扇区。
• 需要改进访问模式，让每次请求真正利用整条 128 B（4 个 32 B 扇区），减少跨更多 Cache line 的开销。

3. DRAM 访问的扇区利用

• DRAM 读取的最小粒度是 64 B（两个 32 B 扇区），当前只用到一半，导致 2 097 154 扇区读取冗余。
• 应避免 ≥64 B 的跨步访问，让每次 DRAM 读都能用到返回的两个扇区，从而提升 DRAM 带宽利用率。

总之，就是在各层缓存和 DRAM 中，都要保证线程束访问既 连续 又 自然对齐，每次事务都能 “恰好” 覆盖所需数据、无多余扇区

6. 对齐但访问地址不是连续的线程ID（add3）

再来看看 add3 的实现：

void __global__ add3(float* x, float* y, float* z){int tid_permuted = threadIdx.x ^ 0x1;int idx = blockIdx.x * blockDim.x + tid_permuted;z[idx]  = x[idx] + y[idx];
}

在 add3 里，虽然我们用

int tid_permuted = threadIdx.x ^ 0x1;
int idx = blockIdx.x * blockDim.x + tid_permuted;

打乱了线程和元素的对应关系，但 访问的元素集合 和 对齐情况 并没有变：每个 warp 仍然要读写同一段从 base+0 到 base+31（共 32 个 float，即 128 B）连续、且恰好以 128 B 对齐的地址空间，只是线程映射被打乱：

• add1 时，线程 0→x[0]、1→x[1]、…、31→x[31]
• add3 时，线程 0→x[1]、1→x[0]、2→x[3]、3→x[2]、…、31→x[30]

因此在物理层面上 add3 依然能合并成 4 个 32 B sector 的事务，和 add1 一样实现了最优的合并访存，带宽利用率仍是 100%

7. warp 中所有线程请求相同地址（add4）

接着看 add4 的实现：

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

在 add4 中，warp 中所有 32 个线程都去读同一个地址 x[warp_idx]（和同一个 y[warp_idx]），这恰好落在一个 对齐 的缓冲行（128 B）或多个 对齐 的 DRAM 传输段（32 B）里，硬件会把这些相同地址的请求 “合并” 成一次 广播（broadcast）事务：

1. DRAM 事务粒度是 32 B

• 对所有线程的请求，只需要发起 1 次 32 B 的 DRAM 事务，就把那 4 B（一个 float）的值读回，然后广播给整个 warp
• 带宽利用率 = 4 B / 32 B = 12.5%

2. L2→L1 Cache line 粒度是 128 B

• DRAM 拉回来的 32 B 事务常常会顺带填充整个 128 B 的缓冲行到 L2/L1
• 这时，只有 4 B 是本次指令实际用到的，带宽利用率 = 4 B / 128 B = 3.125%

因此，add4 虽然所有线程合并到一次最小事务，但只用了非常小的一部分数据，所以效率极低

Nisght Compute 分析结果如下：

相比于 add1 那种每个线程都加载不同地址、造成大量扇区事务的模式，add4 的 Memory Chart 有以下关键差异：

1. 极高的 L1 缓存命中率

• add4 中同一 warp 的 32 条线程都访问完全相同的地址，硬件在 L1 里只需要一次缺失加载，随后 31 次都是 L1 命中 → L1 Hit Rate 飙升到 56.7%（add1 为 0%）。

2. 显著减少的跨层流量

• L1→L2 的传输从 add1 的 ~100 MB（67 MB 读 + 33 MB 写）骤降至 约 11.3 MB（7.5 MB 读 + 3.8 MB 写）；
• 进一步，DRAM 读写流量从近 100 MB 降到 仅 ≈2.43 MB（2.10 MB 读 + 0.33 MB 写）。

3. 更高的 L2 命中率

• L2 缓存命中率上升到 81%（add1 为 33%），配合 L1 的广播机制，使得绝大多数请求都在片上缓存解决，几乎不走 DRAM。

以上变化直接反映了 add4 利用了“广播”+“合并访存”特性：同一地址多重请求被合并并缓存，极大地压缩了内存带宽开销

对应的 Memory Tables 如下图所示：

8. 对齐但非连续访问（add5）

最后看 add5 的实现：

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

add5 中我们将所有线程的 idx 都乘上 4，这也是为什么我们在 add1 实现时只处理 1/4 的数据，但是却开辟了四倍的空间，就是为了防止在测试 add5 的时候出现越界错误

在 add5 里，索引计算改成了

int idx = (blockIdx.x * blockDim.x + threadIdx.x) * 4;

也就是每个线程访问的是

x[base + 0], x[base + 4], x[base + 8], …, x[base + 4·31]

对应的字节地址是（sizeof(float) = 4 B） 是

base_addr + 0·16 B, base_addr + 1·16 B, …, base_addr + 31·16 B

此时：

1. 对齐性

• base_addr 本身是 (blockIdx.x * blockDim.x) * 4 * 4B 且 blockDim.x 是 32 的倍数，因此 base_addr 就是 128 B 对齐的
• 起始地址满足 “自然对齐” 要求

2. 连续性

• 线程访问的地址并不连续，相邻线程间隔 16 B，不是紧贴在一起的 4 B
• 从地址分布上看，这 32 个线程访问点落在 [base_addr, base_addr + 496 B] 这段 500 B 的空间里

3. 内存事务次数

• 全局内存事务的最小单位是 32 B sector，要覆盖 0…499 B，总共需要 $\lceil 500/32\rceil =16$ 个 sector（16 x 32 B = 512 B）才能把这 32 次访问对应的所有 sector 都取回来
• 每个 warp 对 x 的读取会分散到 16 个 sector ⇒ 16 次 32 B 事务
• 同理，对 y 也 16 次，写 z 也会是 16 次—总共 48 次事务

4. 总线利用率

• warp 真正用到的数据量是 32 线程 x 4 B = 128 B
• 但每个数组要搬 16 x 32 B = 512 B
• 利用率 = 128 B / 512 B = 25%

Nisght Compute 分析结果如下：

与 add1 相比，add5 的 Memory Chart 有几个关键区别：

1. 扇区数倍增

• add1 每个 warp 的 128 B 连续区间只跨 4 个 32 B sector；
• add5 因为索引步长 16 B，把访问区间拉长到 512 B，导致每个 load/store 请求要拆成 16 个 32 B sector。

2. 层间流量大幅增加

• L1→L2 总传输量从 add1 的 ≈100 MB（67 MB 读 + 33 MB 写）暴增到 ≈402 MB（268 MB 读 + 134 MB 写）；
• Device Memory 读写流量也分别从 ≈67 MB/32 MB 增至 ≈268 MB/133 MB。

3. 缓存命中率变化很小

• add5 依旧没有 L1 命中（0%），L2 命中率保持在 ≈33%，但因请求扇区数增加，依然有三分之二的流量要下钻到 DRAM，导致总体带宽开销剧增。

Memory Tables 如下所示：

OK，以上就是几种非合并访存的情形分析

9. 总结

最后我们作个简单的总结

写好 CUDA 核函数、最大化合并访存时，主要要注意以下几点：(from ChatGPT)

1. 保证 Warp 访问的起始地址自然对齐

• 通常以 128 B（32-thread * 4B）为自然边界，对齐到 128 B 能让一个 warp 恰好跨整数组块，最优地拆分成 4 个 32 B sector、

2. 让同一 Warp 内的线程访问连续地址

• 线程 threadIdx.x = 0...31 应该映射到物理上相邻的元素（如 x[base + threadIdx.x]），避免跨越多个不相邻的 sector

3. 避免步长多大的稀疏访问

• 如果把索引乘以大于 1 的常数（如 * 4），会把访问分散到更宽的地址范围，导致一次 warp 要发更多 sector 事务、浪费带宽

4. 不要随意给基地址加不对齐的偏移

• 类似 +1 或 ^1 这种打乱对齐的操作，会把紧凑的连续区域 “滑出” sector 边界，增加额外事务

5. 必要时用向量化类型

• 例如用 float4* 加载 / 存储，可一次拿回 16 B（4x4 B），在对齐和连续的前提下，进一步减少指令数

6. 考虑内存访问分支与线程分歧

• 分支内的不同访问模式会破坏统一的合并策略，尽量让全 warp 执行相同的访问模式，例如我们前面讲 reduce 的时候其中的交错寻址就存在线程分歧的问题

7. 关注设备缓存层级策略

• 全局访问默认走 L1→L2→DRAM，保证合并访存也能更好地命中缓存，进一步放大带宽优势

归纳一句话：让一个 warp 地 32 条访问恰好落在一个或几个对齐地连续内存块里，就能最少发起事务、最大化带宽利用

Tips：最后分享一个小技巧，当我们写的 Kernel 比较复杂时不知道哪个地方导致访存没有合并时，Nsight Compute 其实会告诉我们

我们翻到下面的 Source Counters 标签，然后点击对应的源码链接，此时它会跳转到我们的源代码中，并高亮显示导致访存未合并的那行代码，如下图所示：

结语

这篇文章我们跟随 UP 学习了 CUDA 调优中的合并访存，首先我们先了解了相关概念，一个 warp 内的线程对 global memory 进行访问时，如果访问地址是连续且对齐的，那么硬件就可以将这些小访问合并成一次或少数几次大的访问，从而提高带宽利用率、减少延迟

接着我们分析了完美合并访存的核函数代码，并使用 Nsight Compute 进行了验证，然后我们还看了几种访存没合并的情形，并分析了其原因

最后，我们简单做了一个总结，关于写 Kernel 时应注意尽量让同一 warp 内的线程访问连续且对齐的地址空间

OK，以上就是关于 CUDA 调优中合并访存的全部内容了，大家感兴趣的可以多看看 UP 的讲解

下节我们来学习 GPU 中缓存以及访存的具体流程，敬请期待🤗

下载链接

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

参考

• 【CUDA调优指南】合并访存
• CUDA编程学习笔记-03(内存访问)
• NVIDIA CUDA 编程指南
• https://docs.nvidia.com/cuda/cuda-c-programming-guide/
• https://docs.nvidia.com/nsight-compute/NsightCompute