CUDA 内存层次：从访问地址到片上复用

访存优化与数据复用

第五讲已经跑通“分配 → 上传 → kernel → 下载 → 校验”，也看到了传输与缓存怎样影响时间。

向量加法中，线程读取自己的元素。本讲保持结果不变，研究一组线程怎样把数据读得更顺、用得更多。

基础是前四讲的存储层次与复用思想；这次把它们写成可运行的 CUDA 代码。

转置实验的编译与运行

课堂演示使用四个独立实验程序。以下命令在 x64 Native Tools Command Prompt（CMD） 中执行：

cd .\课堂演示\06-memory-hierarchy
code .
nvcc -O3 -std=c++17 -arch=sm_120 ^
  -lineinfo -Xcompiler /O2 06_transpose.cu -o transpose.exe
.\transpose.exe --check-only
.\transpose.exe 4096 4096 20

4096 4096 是输入的行数 M、列数 N，20 是每批重复次数。先看到 PASS，再比较四行 median_ms；稍后解释四个版本。

VS Code 用来阅读和修改源码；改 .cu/.cuh 后重新运行 nvcc，改输入参数则直接再运行 .exe。sm_120 对应课程 RTX 5090 Laptop GPU，其他设备按第五讲方法查询后选择。

实验代码结构与输出说明

PASS 来自与 CPU 参考结果逐项比较；resources 是编译后 kernel 的资源信息。

转置、GEMM 的 benchmark 预热3次、测5批；Stencil 的 compare_kernels 每版预热20次、交替顺序测5批。均显示每批平均每次调用时间的中位数。三者只计 GPU 事件之间的重复 kernel 区间，不计分配、Host/Device 复制和 CPU 校验，也可能包含提交间隙。

Stream 实验另用包含传输的总时间，届时单独说明。

复习：从代码看数据的访问范围

先看同一个 block 中两个线程交换数据。这里用 <<<1, 2>>> 启动，设备数组 a 的前两个数是 10、20：

__global__ void swap_pair(const float* a, float* b) {
    int t = threadIdx.x;
    float mine = a[t];
    __shared__ float tile[2];
    tile[t] = mine;
    __syncthreads();
    b[t] = tile[1 - t];
}

Shared memory 与 L2 缓存的区别

Shared memory 不是 L2。 两者都在 GPU 芯片内，但位置和使用方式不同：

__shared__ float tile[32];
tile[t] = a[i];  // a 指向 global 数组，tile 是 shared 数组

先读取 a[i]：数据可能由 L1 / L2 提供，缓存未命中时才需继续到显存取数；然后写入 shared 中的 tile[t]，供 block 内线程复用。

把输入装入 shared 后反复使用，可以减少重复的 global 读取请求。 Global 是地址空间，不等于每次都访问显存，因此请求减少也不等于显存流量同比减少。

L1 缓存与 Shared memory 的关系

RTX 5090 Laptop 采用这种结构：划给 shared 的容量增大，可用于 L1 的部分减少。色条不代表实际比例。

Global 是地址空间；本讲数组在显存，读取可命中 L1/L2。Shared 由程序显式读写和管理。

已有缓存时，Shared memory 的作用

L1、L2 已经能减少显存访问。Shared 的价值在于让程序明确安排：保留哪些数据、由哪些线程共用、按什么顺序交接。

典型过程是：合作装入 → 同步 → 使用 shared 中的数据；覆盖它之前，还要确认其他线程已用完。

Shared 并不保证比 L1 命中更快。是否值得使用，要比较节省的访存与新增的搬运、同步开销；后面用三个实验观察。

连续访问：32 个线程读取相邻元素

看一个 warp 执行 value = a[lane]：lane 为 0…31，每个线程读一个 4 B 的 float。

把地址按 32 B 一格划分，每格能放 8 个 float；这种小段称为 sector。假设 a 的起始地址按 32 B 对齐：

32 个线程的请求只覆盖 4 格，整格都用上了。 这就是合并访存容易高效的情形；4 格描述地址覆盖，不等于显存一定发生 4 次传输。

跨行访问：每个线程连续，warp 却分散

假设矩阵每行 32 个 float，线程 lane 读取第 lane 行中的第 k 列：

value = a[lane * 32 + k];

竖着看，线程 0 依次读 a[0]、a[1]，确实连续。横着看 k=0，相邻线程却相隔 128 B；完整 warp 的 32 个请求落在 32 个 sector 中。

所以判断合并访存，要看同一轮中相邻线程的地址。下一节的转置，恰好会遇到这种跨行访问。

矩阵转置：从 2×3 到 3×2

转置就是行变成列，列变成行，元素的值不变：

按行存入内存时，A 排成 1、2、3、4、5、6，B 排成 1、4、2、5、3、6。

共享内存的思路：先连续读入一个小块 → 等大家装好 → 换人连续写出。例如读入数字 4 的线程，可以把它交给另一个线程，与数字 1 一起写成 B 的第一行。

朴素矩阵转置：一个线程搬一个元素

把刚才的 2×3 推广为 M×N：A 为 M×N，B 为 N×M，均按行存储；目标是 B[x,y] = A[y,x]。先看 06_transpose.cu 的朴素 kernel：

__global__ void naive(const float* a, float* b, int m, int n) {
    int x = blockIdx.x * 32 + threadIdx.x; // A 的列
    int y = blockIdx.y * 32 + threadIdx.y; // A 的行
    if (x < n && y < m)
        b[size_t(x) * m + y] = a[size_t(y) * n + x];
}
// Host：每块 32×32 个线程，每线程处理一个元素
dim3 block(32, 32), grid((N + 31) / 32, (M + 31) / 32);
naive<<<grid, block>>>(d_a, d_b, M, N);

同一 warp 的 x 连续：读 A 时相邻线程取相邻元素；写 B 时相邻线程跳到不同的行，相隔 M 个 float。

size_t(y) * n + x 就是“行号 × 每行元素数 + 列号”；size_t 让乘法使用较宽的索引类型。

共享内存分块转置的基本原理

课程原创图解 · 点击可查看矢量原图。

线程先合作搬入一个输入 tile，再从 tile 的另一方向取数。写回时，block 的二维坐标也交换。

数学转置由地址映射决定，shared memory 只是让这次映射更适合硬件搬运。

共享内存转置：完整 kernel

仍用 32×32 个线程，每线程装一个、写一个。先完整看 shared32 的核心：

__global__ void shared32(const float* a, float* b, int m, int n) {
    __shared__ float tile[32][32];
    int tx = threadIdx.x, ty = threadIdx.y;
    int x = blockIdx.x * 32 + tx;
    int y = blockIdx.y * 32 + ty;
    if (x < n && y < m) tile[ty][tx] = a[size_t(y) * n + x];
    __syncthreads();
    int out_x = blockIdx.y * 32 + tx;
    int out_y = blockIdx.x * 32 + ty;
    if (out_x < m && out_y < n)
        b[size_t(out_y) * m + out_x] = tile[tx][ty];
}

先按行装入 → 全块同步 → 换人按行写出。源码用 tiled<32> 表示这个版本；<32> 是选定数组行跨度的 C++ 模板参数。改成 tiled<33> 只把 shared 每行长度改为 33，线程数不变。

转置数据写回：两处坐标交换

int out_x = blockIdx.y * 32 + tx; // B 的列：来自 A 的行块
int out_y = blockIdx.x * 32 + ty; // B 的行：来自 A 的列块
if (out_x < m && out_y < n)
    b[size_t(out_y) * m + out_x] = tile[tx][ty];

交换 block 坐标，把整块放到 B 的正确位置；交换 tile 下标，把块内的行变成列。相邻 tx 仍向 B 的连续列写入。

例如 A[1,4]：线程 (tx=4,ty=1) 装入 tile[1][4]，线程 (tx=1,ty=4) 将它写到 B[4,1]。

块内同步与数据可见性

转置的 barrier 位于合作写 shared与从其他线程写的位置读取之间。它同时提供 block 内的会合与相关内存可见性保证。

// 正确结构：谓词只包围有效读写
if (valid_input) tile[...] = a[...];
__syncthreads();
if (valid_output) b[...] = tile[...];

不要把整个结构放入随线程变化的 if (valid)，再让部分参与者跳过同步。设计协作时应明确所有需要的数据由谁提供。

Shared memory 的 32 个 bank

Shared memory 内部分成 32 个可以并行服务的存储分组，称为 bank。一个 bank 可以放很多个数。

看一个从 bank 0 开始的 float s[]：每个 float 占 4 B，相邻元素依次分给相邻 bank，分完 32 个再从头开始。

因此，本例 bank 编号 = 元素下标除以 32 的余数。例如 s[0]、s[32] 在同一个 bank，但它们是两个不同位置。

判断冲突，要看同一 warp 的同一条 shared 访问指令：不同线程的请求落到哪些 bank？请求的是不是同一个位置？

四个线程的读取：无冲突、冲突与广播

同一个 bank、不同位置 → 分开服务；同一个位置 → 读取一次，广播给多个线程。 广播讨论的是读操作。

转置中的 padding：为什么 32 改成 33

转置写回时，同一 warp 的线程沿 shared 的一列读取 tile[tx][0]。先列出线程 0、1、2、3 的地址：

改为 __shared__ float tile[32][33];，只是给每行末尾空出一个位置，让下一行起点错开一个 bank。

完整的 32 个线程也遵循这个规律：跨度 32 时产生 32 路冲突；跨度 33 时分别使用 bank 0…31，没有冲突。

有效数据仍是 32×32，每个 block 多用 32×4 = 128 B。线程数、计算结果和 global 数组的布局都不变。

矩阵转置的编译与性能测试

.\transpose.exe 4096 4096 20
.\transpose.exe 1024 1024 100

四版都采用一线程一个元素，使用 block(32,32)，各自校验正确结果。copy 只作为搬运参考，不能拿复制结果充当转置结果。

RESULT 最后一项是逻辑有效带宽：一个 FP32 读、一个写，共 8MN 字节除以时间；不等于实测显存总线流量。

矩阵转置实测：输入规模与加速比

RTX 5090 Laptop GPU，CUDA 12.8；四版均为 32×32 线程、一线程一个元素。数据已在 GPU，表内为 5 批的 kernel 区间中位数，单位 ms。

• shared 中转有收益：两种规模下，shared32 都比 naive 快约 1.4 倍。
• padding 进一步改善：相对 shared32，小矩阵快约 1.56 倍，大矩阵快约 1.26 倍；不是消除 32 路冲突就快 32 倍。
• 最终对照：padded33 相对 naive 分别快约 2.18 倍、1.73 倍。不同规模的瓶颈可能不同，不能只用 bank 数预测时间。

copy 是参考而非硬件极限；微秒级差异还受提交、频率和批次波动影响。

GEMM 复用：一份输入参与多个输出

回顾一个 2×2 矩阵乘法，让四个线程各算一个输出：

C[0,0] 和 C[1,0] 也都需要 B 的第一列 5、7。独立读取时，不同线程会反复请求相同输入。

让四个线程合作把 A、B 的 8 个数装入 shared，等装好后再各自计算：每个输入供两个输出使用。每线程仍有自己的结果，共享的是输入。

GEMM 的 grid：覆盖整个输出矩阵

constexpr int T = 16; // 输出块边长
dim3 block(T, T);     // 每个 block：256 个线程，各算一个输出
dim3 grid(2, 2);     // C 为 32×32：四个 block 覆盖整个矩阵
gemm_tiled<<<grid, block>>>(dA, dB, dC, 32, 32, K);

一般的 C[M,N] 使用 grid((N+T-1)/T, (M+T-1)/T)，向上取整覆盖边缘。输出块写到不同位置，组成完整的 C；沿 K 的部分和在每个线程内部累加。

朴素 GEMM 中的重复读取

沿用上一页的 grid 分工，计算 C[M,N] = A[M,K] B[K,N]。先看朴素实现：每线程负责一个输出，沿完整 K 逐项累加。

// 06_gemm.cu：T=16，block(T,T)，每线程一个 C[row,col]
int row = blockIdx.y * T + threadIdx.y;
int col = blockIdx.x * T + threadIdx.x;
if (row < M && col < N) {
    float acc = 0;
    for (int k = 0; k < K; ++k)
        acc = fmaf(A[size_t(row)*K+k], B[size_t(k)*N+col], acc);
    C[size_t(row)*N+col] = acc;
}

A[...]、B[...] 每次各取一个数；fmaf(a,b,acc) 计算一次 FP32 融合乘加 a*b+acc。整个 K 循环完成一个输出的点积，相邻输出会反复读取相同输入。

GEMM 输出分块与输入复用

现在放大 grid 中的一个 block：负责一块 T×T 的 C，本例 T=16。输出 C[row,col] 在某个 k 上需要 A[row,k] 和 B[k,col]。

用 TK 表示沿 K 每轮取的项数，本例也取 16。先合作装入 T×TK 个 A + TK×T 个 B，再计算本段对 T×T 个输出的贡献。

输入块在 shared 中共用；每个线程自己的 acc 保留部分和。处理完本段，再装下一段 K。

沿 K 维度的分块累加

课程原创图解 · 点击可查看矢量原图。

T 决定输出块的大小，TK 决定每轮沿 K 累加多少项。 本例 T=TK=16：一个 block 有 256 个线程，每轮装入 A、B 各 256 个 FP32 数，共占 2 KiB shared。

图中只展开一个 block 的计算；其他 block 用同样的方法计算各自的输出块。每线程把 acc 保留在寄存器中，累加完整 K 后写回自己的 C 元素。

GEMM 输入加载与越界补零

constexpr int T = 16;  // 输出块边长
constexpr int TK = 16; // K 方向每段长度
__shared__ float As[T][TK], Bs[TK][T];
int tx = threadIdx.x, ty = threadIdx.y;
int row = blockIdx.y*T + ty, col = blockIdx.x*T + tx;
float acc = 0.0f;
for (int base = 0; base < K; base += TK) {
    As[ty][tx] = (row<M && base+tx<K)
        ? A[size_t(row)*K + base+tx] : 0.0f;
    Bs[ty][tx] = (base+ty<K && col<N)
        ? B[size_t(base+ty)*N + col] : 0.0f;
    __syncthreads();
    // 下一页：计算与再次同步
}

base 是本轮在 K 方向的起点，依次为 0、16、32…。本页每线程各搬一个 A、B 元素的写法要求 T=TK，两者都取16；取不同值时需调整加载分工。越界位置补0，所有线程都参与同步。

GEMM 分块计算与共享缓冲同步

float acc = 0.0f;  // 每线程一份，只初始化一次
for (int base = 0; base < K; base += TK) {
    // 上一页：加载本段 As、Bs，并执行第一次 __syncthreads()
    for (int k = 0; k < TK; ++k)
        acc = fmaf(As[ty][k], Bs[k][tx], acc);
    __syncthreads();  // 本段用完，才能覆盖缓冲
}  // 外层循环结束：已经累加完整 K
if (row < M && col < N)
    C[size_t(row)*N + col] = acc;

k 是本段内的下标 0…TK−1，对应完整 K 方向的 base+k。下一页把输入加载和计算合在一起。

两道 barrier 分别保证：本段数据装好后再读，所有线程用完后再覆盖。

GEMM 分块计算的完整过程

constexpr int T = 16, TK = 16; // 输出块边长、K 分段长度
static_assert(T == TK); // 当前加载方式要求二者相等
__global__ void gemm_tiled(const float* A, const float* B, float* C, int M, int N, int K) {
    __shared__ float As[T][TK], Bs[TK][T];
    int tx = threadIdx.x, ty = threadIdx.y;
    int row = blockIdx.y*T + ty, col = blockIdx.x*T + tx;
    float acc = 0; // 每线程一个输出的累加器，只初始化一次
    for (int base = 0; base < K; base += TK) {
        As[ty][tx] = (row<M && base+tx<K)
            ? A[size_t(row)*K + base+tx] : 0.0f;
        Bs[ty][tx] = (base+ty<K && col<N)
            ? B[size_t(base+ty)*N + col] : 0.0f;
        __syncthreads(); // 本段输入全部装好
        for (int k = 0; k < TK; ++k)
            acc = fmaf(As[ty][k], Bs[k][tx], acc);
        __syncthreads(); // 所有线程用完，才能覆盖为下一段
    }
    if (row < M && col < N) C[size_t(row)*N + col] = acc;
}

dim3 block(T, T), grid((N+T-1)/T, (M+T-1)/T);
gemm_tiled<<<grid, block>>>(dA, dB, dC, M, N, K);

每段：加载 → 同步 → 累加 → 同步；全部 K 段结束后，只写回一次。

GEMM 动画：从全局到块内计算

GEMM 尾部处理：K=18 的示例

每个输出的点积需要完整 K 下标 0…17。TK=16 时共两轮：base=0 覆盖 0…15，base=16 覆盖 16…31。

最后一轮中，段内 k=0、1 对应有效输入 16、17；k=2…15 对应的输入已补零。每个输出仍执行 TK 次乘加，额外的零乘加不改变精确数学结果。

M、N 的尾部也要区分：无效输出线程不能写 global，却可能仍参与加载别人需要的 shared 元素与两道 barrier。计算边界、加载边界、同步边界必须分别判断。

分块 GEMM 的访存请求量分析

只看一个 T×T 输出块、K 方向的 TK 项，本例 T=TK=16，不含尾部：

每个装入 shared 的输入值参与 16 个输出的计算；算术工作量不变。新增代价是 shared 读写、两次同步和 shared 容量。

减少 16 倍的是源码层的输入读取请求，不是保证快 16 倍。 朴素版已有缓存与同字访问复用；只有测量才能判断净收益。

GEMM 的编译与性能测试

nvcc -O3 -std=c++17 -arch=sm_120 ^
  -lineinfo -Xcompiler /O2 06_gemm.cu -o gemm.exe
.\gemm.exe --check-only
.\gemm.exe 512 512 512 20
.\gemm.exe 1024 1024 1024 20

参数是 M N K 每批次数。CPU FP64 参考与 GPU FP32 结果按容差逐项校验；CPU 校验不计入事件时间。

两个命令只改变矩阵尺寸。下一页对照结果，观察shared复用是否带来收益。

GEMM 实测：数据复用与加速比

同一 RTX 5090 Laptop GPU、CUDA 12.8，数据已在 GPU；表内为5批 kernel 区间平均时间的中位数，每批20次。这次简单分块有收益，但远不到16倍，也不代表高性能 GEMM 的上限。

Stencil：相邻输出共享哪些输入

设输入 x = [1, 2, 3, 4, 5, 6, 7]，下标从 0 开始。每个输出取“自己与左右各两个邻居”的平均值：

两个线程共用了 2、3、4、5 这四个数；分别读取要发出 10 次请求，但实际只涉及 6 个不同的输入。

共享内存的思路：把共同需要的一段输入装好 → 同步 → 每个线程取自己的 5 个数求平均。

推广到更长的数组：若一组线程负责中间连续 8 个输出，就需要 左侧 2 个 + 中间 8 个 + 右侧 2 个 = 12 个输入。两侧额外的邻居区域称为 halo。

一维 Stencil：半径为 2 的邻域平均

把刚才的邻域平均写成一般形式：输入 x 长度为 N，输出 y 长度相同。数组边缘缺少的邻居视为 0，分母仍为 5；x、y 不重叠。

$y_i=\frac{x_{i-2}+x_{i-1}+x_i+x_{i+1}+x_{i+2}}{5}.$

float sum = 0;
for (int d = -2; d <= 2; ++d) {
    int j = i + d;
    if (j >= 0 && j < n) sum += x[j];
}
y[i] = sum * 0.2f; // 前提 i<n，索引规模在 int 范围内

仍用上一页的 x=[1,2,3,4,5,6,7]，y[0]=(0+0+1+2+3)/5=1.2；边界也除以5。

前面用转置重排访问、用 GEMM 复用输入；现在把两种判断迁移到第三讲的邻域计算。一个输出需要中心及邻居，因此 tile 除了主体还要带 halo。

Stencil 输入分块与 halo 区域

课程原创图解 · 点击可查看矢量原图。

一个 block 负责 B 个连续输出，需要 B+2R 个输入；其中左右各 R 个属于 halo。

R=2、B=8：输出索引 8…15，需要输入 6…17。左右 halo 不属于本 block 的输出，但必须有人加载。

naive 对每个输出发出 5 次输入读取；shared 版本把重叠窗口先装一次，再重复使用。多个线程读同一输入无需互斥，协作装入后需要明确完成边界。

Stencil 索引：全局数组与 shared 位置

kernel 接收输入指针 x、输出指针 y 和数组长度 n，两数组都在 GPU 上。B=256 是每块的线程数，也是每块负责的输出个数；以下分步代码先用邻域半径 R=2。

沿用前图的 B=8、base=8：tile[0] 放 x[6]，tile[2] 放 x[8]，tile[11] 放 x[17]。减 R，是为左侧 halo 留出位置。

线程 t=0 负责 y[8]，计算时需要 x[6…10]，也就是 tile[0…4]。加载位置 j 与输出位置 i 各有用途。

Stencil 的协作加载与块内同步

constexpr int B = 256, R = 2;
__shared__ float tile[B + 2*R];
int t = int(threadIdx.x);          // 块内线程编号
int base = int(blockIdx.x) * B;    // 本块的输出起点
for (int j = t; j < B + 2*R; j += B) {
    int g = base + j - R;          // shared 位置 → 全局输入
    tile[j] = (g >= 0 && g < n) ? x[g] : 0.0f;
}
__syncthreads();

256 个线程装入 260 个位置：j += B 表示一个线程装完一处，再向后跨 256 个位置。

条件 ? x[g] : 0.0f：下标有效就读输入，否则补 0。装入后全体执行 __syncthreads()，再开始读 shared；最后一块里没有有效输出的线程也要参与加载。

Stencil 输出计算与边界处理

int i = base + t; // 当前线程负责的全局输出
if (i < n) {
    float sum = 0;
    for (int d = 0; d < 2*R + 1; ++d)
        sum += tile[t + d];
    y[i] = sum * 0.2f; // R=2，共5项，乘0.2即除以5
}

这里 d=0…4 是窗口内的偏移。由上一页的映射可知：tile[t+d] 对应 x[i-R+d]，因此正好读到 x[i−2]、x[i−1]、x[i]、x[i+1]、x[i+2]。

例如线程 t=0 读 tile[0…4]，线程 t=1 读 tile[1…5]，其中4个输入共用。边界补零已经在加载时完成，这里仍固定除以5。

先全体加载并同步，再用 i<n 判断谁写输出。 本 kernel 只使用这一段输入，之后不再覆盖 tile，因此不需要 GEMM 中“覆盖前”的第二道 barrier。

扩大邻域：让协作加载带来更多复用

5点窗口为什么可能更慢？ 朴素版的访问连续、相邻窗口高度重叠，L1/L2 已有机会复用输入。shared 版还需写入 tile、同步一次、再从 tile 读取；窗口小时，这些额外步骤可能抵消收益。

保留刚才的加载方式，把实验改为 R=32：左右各32个邻居，共65项的滑动平均。它对应更宽的平滑窗口，每个输出仍由一个线程计算。

输入只多装60个，服务的邻域读取却变为原来的13倍。加载时检查边界、补零一次，后续在 shared 内反复使用。 代码将分母改为 2*R+1，其余分工不变。

以上是程序请求数，不是实测显存流量或加速倍数；接下来在每个半径内公平比较 naive 与 shared。

Stencil 共享内存计算的完整过程

constexpr int B = 256, R = 32; // 本页采用65点窗口
__global__ void stencil_shared(const float* x, float* y, int n) {
    __shared__ float tile[B + 2*R];
    int t = int(threadIdx.x), base = int(blockIdx.x) * B;
    for (int j = t; j < B + 2*R; j += B) {
        int g = base + j - R;
        tile[j] = (g >= 0 && g < n) ? x[g] : 0.0f;
    }
    __syncthreads(); // 本块需要的输入已全部装好
    int i = base + t;
    if (i < n) {
        float sum = 0;
        for (int d = 0; d < 2*R + 1; ++d) sum += tile[t + d];
        y[i] = sum / float(2*R + 1);
    }
}

dX、dY 是已分配的设备数组，输入已上传至 dX；本例 0≤n≤2²⁴：

if (n > 0) {
    int blocks = (n + B - 1) / B; // 向上取整，覆盖全部输出
    stencil_shared<<<blocks, B>>>(dX, dY, n);
    CUDA_CHECK(cudaGetLastError());
    CUDA_CHECK(cudaDeviceSynchronize());
}

定位输出块 → 合作装入320个输入 → 同步 → 各算65项平均 → 写回。

Stencil 实验：编译与对比运行

nvcc -O3 -std=c++17 -arch=sm_120 ^
  -lineinfo -Xcompiler /O2 06_stencil.cu -o stencil.exe
.\stencil.exe --check-only
.\stencil.exe 4194304 100
.\stencil.exe 16777216 100

默认编译 R=32。输出依次显示 N、R、window、两版校验结果，以及 naive / shared 时间和 Speedup；Speedup > 1 表示 shared 更快。

若想复现刚才的5点窗口，用同一源码再编译一次：

nvcc -O3 -std=c++17 -arch=sm_120 -lineinfo -Xcompiler /O2 ^
  -DSTENCIL_RADIUS=2 06_stencil.cu -o stencil_r2.exe
.\stencil_r2.exe 4194304 100

-D 在编译时指定半径；两个 .exe 各自比较同一窗口下的两种实现。每版预热20次，交替先后顺序测5批，每批100次；只计 kernel，不计分配、复制和 CPU 校验。

Stencil 实测：宽窗口下 shared 获得收益

在 RTX 5090 Laptop GPU、CUDA 12.8 上，使用相同输入、block=256及编译选项。每个半径内的 naive 与 shared 结果均通过 CPU 参考校验。

时间取3轮独立运行的中位数。R=32各轮均更快，加速比范围为 1.23–1.42×。

小窗口：额外加载与同步不一定划算。宽窗口：更多线程重复使用同一段输入，协作加载的成本被摊薄，本次约快1.4倍。

global 输入请求降到1/52，也不等于52倍加速：朴素版也有缓存，shared 版仍要执行65项读取与加法，并承担加载、同步等开销。

CUDA stream：任务队列与执行顺序

前面的 shared 解决 一个 kernel 内的数据协作。现在回到第五讲的“上传 → kernel → 下载”，观察 多个任务之间怎样安排顺序。

Stream 是提交 CUDA 操作的有序队列。 同一 stream 中，后面的操作等待前面的操作完成；CPU 可以先把它们依次提交出去。

CPU 提交顺序：上传 A，计算 A，下载 A
GPU 队列顺序：[上传 A] → [计算 A] → [下载 A]

以前写 kernel<<<grid, block>>>，用的就是默认 stream，并不是没有 stream。

多个 stream 可以表达互相独立的工作；它们共享 GPU 资源，不保证同时运行，也不各自独占一个 SM。

计算示例：叠加多项正弦信号

对每个输入相位 x[i]，计算 sin(x[i]) + sin(2x[i])/2 + … + sin(terms·x[i])/terms。例如 terms=2 时，只加前两项。

__global__ void synthesize(const float* x, float* y, int n, int terms) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i >= n) return;
    float sum = 0.0f;
    for (int k = 1; k <= terms; ++k)
        sum += sinf(k * x[i]) / k;
    y[i] = sum;
}

terms 是叠加项数，本例取 64；sinf 是单精度正弦函数，输入角度用弧度。每个线程仍只负责一个 y[i]。

不同输出互不依赖，可以把数组分段。 每个采样点需要累加64项正弦值，为计算与传输的重叠提供了机会。

异步提交：CPU 可以先做独立的工作

synthesize<<<grid, block>>>(d_in, d_out, n, terms);
CUDA_CHECK(cudaGetLastError()); // 检查提交错误，不等待计算结束
prepare_next_input();          // 自编 CPU 函数示意，与当前 GPU 缓冲无关
CUDA_CHECK(cudaDeviceSynchronize());
// 等 GPU 完成后，才继续依赖它的步骤

异步说的是“调用返回时，任务可能还没完成”。 CPU 可以在 GPU 计算时准备下一份输入，不必启动多个 stream 才能做到。

若 CPU 做的是读取本次 GPU 结果，就必须等结果可用。若马上调用同步函数，CPU 仍会在这里等待。

“CPU 没在等”与“两个 GPU kernel 同时执行”是两个问题；后者需要独立 stream 和足够的执行资源。

Pinned 主机内存与异步复制

要让 Host↔Device 传输有机会与计算重叠，先准备 pinned 主机内存。h_in 放输入相位，h_out 接收合成结果：

float *h_in, *h_out;
size_t total_bytes = size_t(n) * sizeof(float);
CUDA_CHECK(cudaMallocHost(&h_in, total_bytes));
CUDA_CHECK(cudaMallocHost(&h_out, total_bytes));
// 初始化输入，提交任务，等待完成；然后释放
CUDA_CHECK(cudaFreeHost(h_in));
CUDA_CHECK(cudaFreeHost(h_out));

Pinned 表示这些 CPU 内存页被固定，便于设备直接传输；它没有把数组变成 GPU 显存。

cudaMemcpyAsync(..., s) 指定复制队列。普通 malloc / std::vector 内存可能触发暂存或等待，不能仅凭 Async 名字断定已实现重叠。

数据分段与缓冲区的对应关系

每段（chunk）包含 chunk_size 个元素。A、B、C 三段的总长度 n=3×chunk_size。

例如 chunk_size=4，n=12：

A：c=0，offset=0 → 元素[0…3]
B：c=1，offset=4 → 元素[4…7]
C：c=2，offset=8 → 元素[8…11]

d_in + offset 是本段起点；kernel 收到这个指针后，仍用局部下标 0…chunk_size−1。

同一 stream：上传、计算与下载

先看 A段的完整任务链。下面四个数组均已分配，bytes=chunk_size×sizeof(float)：

cudaStream_t s;
CUDA_CHECK(cudaStreamCreateWithFlags(&s, cudaStreamNonBlocking));
CUDA_CHECK(cudaMemcpyAsync(d_in, h_in, bytes, cudaMemcpyHostToDevice, s));
synthesize<<<(chunk_size + 255) / 256, 256, 0, s>>>(
    d_in, d_out, chunk_size, terms);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaMemcpyAsync(h_out, d_out, bytes, cudaMemcpyDeviceToHost, s));
CUDA_CHECK(cudaStreamSynchronize(s));
CUDA_CHECK(cudaStreamDestroy(s));

启动配置第三项 0 是动态 shared 字节数，第四项 s 才是 stream。上传和 kernel 之间不用插入 CPU 等待；同一队列已保证先后。

最后同步后，CPU 才能读 h_out。NonBlocking 表示这条队列不参与 legacy 默认 stream 的隐式同步。

不同数据段的传输与计算重叠

每段内部仍然是 上传 → 计算 → 下载。不同段相互独立，例如计算A时可以上传B，计算B时可以下载A。

不是把同一段的三个阶段分给三条 stream。 同一段仍有先后依赖；这里每条队列负责一段的完整任务链。

这是可能的执行安排，不是本次测量时间线；能否重叠及收益大小取决于设备、任务规模和提交顺序。

一条与三条 stream：只改变任务归属

streams 是存放队列句柄的数组；stream_count 是本次使用的队列数。用 streams[c % stream_count] 给第c段选择队列，% 表示取余：

CPU 的提交顺序采用三个循环：先上传A/B/C，再计算A/B/C，最后下载A/B/C。每段的三项操作始终选择同一条 stream，依赖不会丢失。

提交完不等于执行完。 CPU开始提交计算时，GPU仍可能在上传其他段。三个循环之间不插入同步，全部提交后再等待。

两种配置的总数据量、计算项数、分段数和复制次数完全相同：单stream串行处理三段，三stream各负责一段。

完整提交过程：三个循环与最后的等待

run_pipeline 是自编的 Host 函数，下面展示全部实现；数组与 stream 由调用方创建。

void run_pipeline(const float* h_in, float* h_out, float* d_in, float* d_out,
                  int chunk_size, int terms, cudaStream_t* streams, int stream_count) {
    size_t bytes = size_t(chunk_size) * sizeof(float);
    for (int c = 0; c < 3; ++c) { // Submit all uploads first.
        int offset = c * chunk_size;
        CUDA_CHECK(cudaMemcpyAsync(d_in + offset, h_in + offset, bytes,
            cudaMemcpyHostToDevice, streams[c % stream_count]));
    }
    for (int c = 0; c < 3; ++c) { // Then submit all kernels.
        int offset = c * chunk_size;
        synthesize<<<(chunk_size + 255) / 256, 256, 0,
            streams[c % stream_count]>>>(d_in + offset, d_out + offset, chunk_size, terms);
        CUDA_CHECK(cudaGetLastError());
    }
    for (int c = 0; c < 3; ++c) { // Finally submit all downloads.
        int offset = c * chunk_size;
        CUDA_CHECK(cudaMemcpyAsync(h_out + offset, d_out + offset, bytes,
            cudaMemcpyDeviceToHost, streams[c % stream_count]));
    }
    for (int s = 0; s < stream_count; ++s)
        CUDA_CHECK(cudaStreamSynchronize(streams[s]));
}

三条队列分别得到 A、B、C 的完整任务链。所有任务已提交，先等队列0不会阻止另外两条继续执行；函数返回时，三个输出段都已回到 h_out。

完整调用：准备、执行、读取与释放

文件开头写下方的 #include，接着放已展示的 synthesize、run_pipeline，最后放 main，即为三 stream 最小程序：

#include "lab06_utils.cuh"
// 此前已定义 synthesize 和 run_pipeline
int main() {
    const int chunk_size = 1 << 20, terms = 64;
    const int n = 3 * chunk_size;
    size_t total_bytes = size_t(n) * sizeof(float);
    float *h_in, *h_out, *d_in, *d_out;
    CUDA_CHECK(cudaMallocHost(&h_in, total_bytes));
    CUDA_CHECK(cudaMallocHost(&h_out, total_bytes));
    CUDA_CHECK(cudaMalloc(&d_in, total_bytes));
    CUDA_CHECK(cudaMalloc(&d_out, total_bytes));
    for (int i = 0; i < n; ++i)
        h_in[i] = float(i % 1024) * 0.006f;
    cudaStream_t streams[3];
    for (int s = 0; s < 3; ++s)
        CUDA_CHECK(cudaStreamCreateWithFlags(&streams[s], cudaStreamNonBlocking));
    run_pipeline(h_in, h_out, d_in, d_out,
                 chunk_size, terms, streams, 3);
    std::printf("y[1] = %.6f\n", h_out[1]); // 全部下载完成后读取
    for (int s = 0; s < 3; ++s) CUDA_CHECK(cudaStreamDestroy(streams[s]));
    CUDA_CHECK(cudaFree(d_in)); CUDA_CHECK(cudaFree(d_out));
    CUDA_CHECK(cudaFreeHost(h_in)); CUDA_CHECK(cudaFreeHost(h_out));
}

调用末尾的 3改成1，就得到单stream版本。课堂演示程序06_stream.cu在同一流程外增加了数值校验和公平计时。

使用 event 建立跨 stream 依赖

前面的三段互不依赖。如果 s1 必须使用 s0 的结果，就需要显式建立联系。下面另用“加1”小例子说明依赖，不参与后面的性能比较：

__global__ void add_one(const float* x, float* y, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) y[i] = x[i] + 1.0f;
}
cudaEvent_t ready;
CUDA_CHECK(cudaEventCreateWithFlags(&ready, cudaEventDisableTiming));
add_one<<<grid, block, 0, s0>>>(d_in, d_tmp, n);
CUDA_CHECK(cudaEventRecord(ready, s0));
CUDA_CHECK(cudaStreamWaitEvent(s1, ready, 0));
add_one<<<grid, block, 0, s1>>>(d_tmp, d_out, n);
CUDA_CHECK(cudaStreamSynchronize(s1));
CUDA_CHECK(cudaEventDestroy(ready));

ready 是 s0 的完成标记。StreamWaitEvent让s1等，CPU仍可继续提交；StreamSynchronize让CPU等。 两次加1后得到输入加2。

这里 d_in 已上传，d_tmp / d_out 已分配，s0 / s1 已创建。仅CPU先提交s0再提交s1，并不能保证跨队列完成顺序。

Stream 实验：编译与运行

沿用课堂演示目录和 CMD 开发终端，重新编译06_stream.cu：

nvcc -O3 -std=c++17 -arch=sm_120 ^
  -lineinfo -Xcompiler /O2 06_stream.cu -o stream.exe
.\stream.exe --check-only
.\stream.exe 1048576 64
.\stream.exe 4194304 64

第一个参数是 每段长度 chunk_size，总元素数是它的3倍；第二个参数是 正弦叠加项数 terms。

程序自动比较1条与3条 stream。先看两行 PASS，再看包含传输的 total 和相对单stream的 speedup。末行 event dependency 只验证输入加2。

每种配置预热5次，交替先后顺序测9次，取中位数。CPU计时包含 上传、计算、下载、提交和完成等待，排除分配、创建队列与数值校验。

Stream 实测：让传输与较长计算重叠

RTX 5090 Laptop GPU、CUDA12.8，固定 3段数据、每点64项正弦叠加。各配置只改变stream归属，均通过CPU参考校验。

表中为3轮独立运行的中位数；三条stream在各轮均更快，加速范围 1.47–1.55×。

三段数据相互独立，多条 stream 可以让一段的计算与另一段的传输重叠，从而缩短总时间。

三条stream不等于三倍速度：它们仍共享计算和传输资源，还有开始与收尾阶段。继续增加计算量或队列数，也不保证继续变快。

本讲小结

• 先看地址：合并访存看同一次指令中，一个 warp 请求的地址集合。
• 再组织协作：转置用 shared 重排，GEMM 用 shared 复用；同步保护装入、读取与覆盖。
• 最后看实测：请求减少、bank 冲突消失、stream 增多，都不能单独决定加速比。
• 任务之间看依赖：同一 stream 保序；独立任务才有机会重叠，CPU 读取结果前仍需等待。

下一讲研究另一种协作：多个线程的结果怎样合成一个和、一个前缀位置，或一张直方图。