CUDA tiled SGEMM 入门：共享内存分块、同步屏障与越界补零

Tiled SGEMM 第一版的核心思路很直接：每个 CUDA block 认领 C 的一个小方块 BM×BN，然后沿着 K 方向一片片切。每一轮把 A 的 BM×BK 竖条、B 的 BK×BN 横条搬进 shared memory；全 block 同步后，每个线程从 smem 取数做乘加，累加到自己的寄存器；K 循环结束后，把累加结果写回 C。

今天只求对拍 PASS，不调性能。第一版通常只有 naive 的 2~4 倍，这很正常。

1. 先认识尺寸和坐标

矩阵乘法：

tile 尺寸：

名字	含义
BM	一个 block 负责的 C 子块行数
BN	一个 block 负责的 C 子块列数
BK	每轮沿 K 方向搬多少

坐标惯用法：

blockIdx.x  -> C 的 N 方向，列方向
blockIdx.y  -> C 的 M 方向，行方向

threadIdx.x -> block 内列方向
threadIdx.y -> block 内行方向

C 被切成网格：

C：M 行 × N 列

          blockIdx.x / N 方向
       0        1        2
   ┌────────┬────────┬────────┐
 0 │ B00    │ B01    │ B02    │
   │ BM×BN  │        │        │
   ├────────┼────────┼────────┤
 1 │ B10    │ B11    │ B12    │
   │        │        │        │
   ├────────┼────────┼────────┤
 2 │ B20    │ B21    │ B22    │
   └────────┴────────┴────────┘
blockIdx.y / M 方向

每个 block 只算其中一个小方块。block 内每个线程算这个小方块里的一个或几个元素。

2. tiled 到底在“tile”什么？

K 循环每轮只取一小段 K：

A：M×K                          B：K×N
┌──────────────┐                ┌──────────────┐
│              │                │              │
│   A 竖条     │  BM×BK         │   B 横条     │  BK×BN
│              │                │              │
└──────────────┘                └──────────────┘
         ×                               =
                 C 子块：BM×BN
        ┌──────────────┐
        │   C_tile     │
        │  累加结果     │
        └──────────────┘

对每个 K 切片：

C_tile += A_tile × B_tile

注意：A 是 BM×BK 的“竖条”，B 是 BK×BN 的“横条”。它们被搬进 shared memory：

shared memory 里每个 block 一份：
sA[BM][BK]     sB[BK][BN]
┌─────────┐    ┌─────────┐
│ A 小块  │    │ B 小块  │
└─────────┘    └─────────┘

shared memory 快，但是每个 block 私有，block 之间不共享。

3. 一个 block 内部的流程

用“搬运工 / 计算工 / 围栏”来理解：

所有线程合作搬 A 竖条、B 横条进 smem
                │
                ▼
        【屏障 1：搬完才准算】
                │
                ▼
每个线程从 smem 取数，做乘加，累加到自己的 acc
                │
                ▼
        【屏障 2：算完才准搬下一轮】
                │
                ▼
下一轮 K 切片，覆盖 smem

对应骨架：

__global__ void sgemm_tiled_v1(...)
{
    __shared__ float sA[BM][BK];
    __shared__ float sB[BK][BN];

    int row = blockIdx.y * BM + threadIdx.y;
    int col = blockIdx.x * BN + threadIdx.x;

    float acc = 0.0f;

    for (int kk = 0; kk < K; kk += BK) {

        // TODO：所有线程合作搬 A 竖条 / B 横条进 smem
        // 搬运任务可以用 if 挑人；越界要补 0
        // 但下面这个屏障不能放进 if 里

        __syncthreads();   // 屏障 1

        // TODO：从 smem 读数，累加到 acc
        // for k in 0..BK-1:
        //     acc += sA[threadIdx.y][k] * sB[k][threadIdx.x]

        __syncthreads();   // 屏障 2
    }

    // 写回 C，注意边界 guard
    if (row < M && col < N) {
        C[row * N + col] = acc;
    }
}

这是一个“一个线程算一个 C 元素”的简化示意。实际作业里如果 BM/BN 和 blockDim 不完全一样，就让一个线程算多个元素。

4. 搬运的索引怎么写？

常用写法是：把 BM×BK 或 BK×BN 拉平成一维任务，每个线程按自己的 tid 去领任务。

A 竖条搬运示意：

int tid = threadIdx.y * blockDim.x + threadIdx.x;
int nthreads = blockDim.x * blockDim.y;

for (int i = tid; i < BM * BK; i += nthreads) {
    int r = i / BK;          // A 小块内的行
    int c = i % BK;          // A 小块内的列，也就是 K 方向偏移

    int gr = blockIdx.y * BM + r;  // A 的全局行
    int gc = kk + c;               // A 的全局列

    sA[r][c] = (gr < M && gc < K) ? A[gr * K + gc] : 0.0f;
}

B 横条搬运示意：

for (int i = tid; i < BK * BN; i += nthreads) {
    int r = i / BN;          // B 小块内的行，也就是 K 方向偏移
    int c = i % BN;          // B 小块内的列

    int gr = kk + r;               // B 的全局行
    int gc = blockIdx.x * BN + c;  // B 的全局列

    sB[r][c] = (gr < K && gc < N) ? B[gr * N + gc] : 0.0f;
}

累加示意：

for (int k = 0; k < BK; ++k) {
    acc += sA[threadIdx.y][k] * sB[k][threadIdx.x];
}

这三段是核心作业：索引算式 + 搬运 guard + 累加逻辑。

5. 两道 __syncthreads() 分别防什么？

时间线图：

第 r 轮：

load sA/sB(r) ──► 【屏障1】 ──► compute(r) ──► 【屏障2】 ──► load sA/sB(r+1)

屏障 1：防“写后读”竞态

没有屏障1：

快线程：load 完自己的部分 ──► 立刻开始 compute
慢线程：还没 load 完
                         ↑
快线程读到还没写好的 smem / 垃圾值

所以屏障 1 保证：所有线程把本轮 smem 搬完，大家才一起开始读。

屏障 2：防“下一轮搬运覆盖本轮还没读完的 smem”

这是今天口述检验点。

没有屏障2：

快线程：compute(r) 完 ──► load(r+1)，开始覆盖 smem
慢线程：还在 compute(r)，仍在读 smem 旧值
                              ↑
                    读到被覆盖后的新值 / 半新半旧 -> 数错

所以屏障 2 保证：所有线程都把本轮 smem 读完，才允许下一轮 load 覆盖它。

一句话答案：

第二道 __syncthreads() 防的是：快线程提前进入下一轮 K 切片，把 shared memory 覆盖掉，而慢线程还在读上一轮 smem 的竞态，也就是跨轮次的 WAR 竞态。

两道屏障一个都不能少，位置也不能挪。

6. 条件分支里的 __syncthreads() = 死锁

__syncthreads() 的铁律：

block 内所有线程必须全部到达屏障，计数器凑齐才放行。

错误写法：

if (tid < load_count) {
    // 搬运
    __syncthreads();   // 错！部分线程永远到不了
}

图示：

线程 0~99：到达屏障 ──┐
线程 100~255：被 if 挡住，永远不到 ──┐
                                      ▼
                             屏障计数器永远凑不齐
                                      ▼
                                 卡死 / 未定义

正确写法是两句话，不要揉成一句：

if (tid < load_count) {
    // 搬运任务：可以挑人干
}

__syncthreads();   // 屏障：全员无条件通过

如果 kernel 出现“跑着跑着卡住、nvidia-smi 里 GPU 利用率 100% 但无进展”，第一怀疑对象就是它。全文搜索 __syncthreads，逐个确认：所有线程是否无条件执行到？

7. 三处边界，分开处理

这一节是今天最容易“不崩不报错但数错”的地方。先把越界到底是什么、为什么会越界讲清楚，再分别看三处边界。

7.1 越界到底是什么？

矩阵的有效范围是：

A：行 0..M-1，列 0..K-1
B：行 0..K-1，列 0..N-1
C：行 0..M-1，列 0..N-1

所谓“越界”，不是“CUDA 一定报错崩溃”，而是：你根据 blockIdx / threadIdx / kk 算出来的全局坐标，已经超出了矩阵真正有效的范围，但代码仍然拿这个坐标去读 A/B 或写 C。

越界分两种：

类型	含义	表现
物理越界	地址真的超出 cudaMalloc 分配范围	段错误 / 非法内存访问
逻辑越界	地址还在大缓冲区里，但读到/写到别的行、别的列	不崩不报错，只是数错

逻辑越界更阴险。因为矩阵是行主序拉平的：

A[gr][gc] 对应一维下标：gr * K + gc

如果 gc == K：

A[gr][K] -> gr*K + K = (gr+1)*K + 0
          -> 实际是 A[gr+1][0]，下一行开头

图示：

A 行主序：

行 gr：   [ ... K 个有效 ... ] [ gc=K 越界 ]
                               ↓
行 gr+1： [ 下一行第 0 个元素 ... ]

gc >= K 时，不是简单多读一个，而是跳到下一行去读。

所以读越界不一定崩，只是悄悄读错数据。只有最后一行再往后越，才容易变成物理越界。

7.2 为什么一定会越界？因为 tile 是固定窗口，grid 只能向上取整

tile 尺寸是固定的：BM×BN，K 方向固定切 BK。但矩阵维度不一定刚好是 tile 的整数倍。

所以 grid 必须向上取整：

gridDim.y = ceil(M / BM)
gridDim.x = ceil(N / BN)

ceil 意味着：最后一个 block 会伸出矩阵边界。

举例：

M = 1000，BM = 16
ceil(1000 / 16) = 63 个 block
63 × 16 = 1008 行

最后一个 block（blockIdx.y = 62）的起始行：

62 × 16 = 992

它覆盖的行：

992, 993, ..., 999, 1000, 1001, ..., 1007
└──── 有效 8 行 ────┘ └──── 越界 8 行 ────┘

图示：

M = 1000，BM = 16

行 0    ┌─────────────┐
        │ block 0     │ 行 0..15
行 16   ├─────────────┤
        │ block 1     │ 行 16..31
        ├─────────────┤
        │ ...         │
行 992  ├─────────────┤
        │ block 62    │ 行 992..1007
行 1000 │─────────────│ ← 矩阵到此结束
        │ 越界 1000..1007 │ ← 矩阵里没有这些行
        └─────────────┘

在 blockIdx.y = 62 的 block 里，threadIdx.y = 8..15 的线程会算出：

row = 62*16 + 8  = 1000  ← 越界
row = 62*16 + 9  = 1001  ← 越界
...
row = 62*16 + 15 = 1007  ← 越界

列方向同理：N = 1000, BN = 16 时，最后一个 block 的 col = 1000..1007 也越界。

K 方向同理：K = 1000, BK = 16 时，最后一轮 kk = 992，k 覆盖 992..1007，其中 1000..1007 越界。

一句话：不是代码写错了才越界，而是“固定 tile + 向上取整”天然会让最后一块伸出边界。

7.3 边界 1：C 子块出界，写越界

C 的写回：

int row = blockIdx.y * BM + threadIdx.y;
int col = blockIdx.x * BN + threadIdx.x;

C[row * N + col] = acc;

如果 row >= M 或 col >= N，这个线程根本不该写 C。

图示：

C 的最后一个 block：

有效区域（row<M, col<N）：
┌──────────────┐
│ 写 acc        │
│ 写 acc        │
│ 写 acc        │
└──────────────┘
越界区域（row>=M 或 col>=N）：
┌──────────────┐
│ 不写          │ ← if 挡住
│ 不写          │
└──────────────┘

写回时 guard：

if (row < M && col < N) {
    C[row * N + col] = acc;
}

逐项读：

条件	含义
row < M	这一行在 C 里真实存在
col < N	这一列在 C 里真实存在
&&	两个都成立才写

核心理解：

row/col 是“我算出来的坐标”，不一定真实存在；if 是在问“C 里真的有这个格子吗”，有才写。

如果不加 if，越界线程会执行：

C[row * N + col] = acc;

比如 row = 1000, N = 1000, col = 5：

下标 = 1000 * 1000 + 5 = 1000005
C 合法下标：0 .. 999999
这次写入：1000005  -> 越界

后果可能是段错误，也可能写到别的内存上，把别的东西改坏。

为什么 C 写回越界“直接跳过”不影响？

因为 C 的越界位置本来就不存在：

C 真实范围：row < M, col < N
越界位置：row >= M 或 col >= N

这些位置不是“C 里应该写 0 的格子”，而是“C 里根本没有的格子”。

所以：

有效格子 -> 写 acc
无效格子 -> 什么都不做

有效 C 元素只由有效线程写，越界线程不写不会影响它们。就像往作业本上写答案：本子只有 1000 行，你算到第 1001 行，直接不写就行，不需要在第 1001 行写 0。

7.4 边界 2：K 切片搬运出界，读越界

A 竖条是 BM×BK：

A 小块的行 -> M 方向
A 小块的列 -> K 方向

B 横条是 BK×BN：

B 小块的行 -> K 方向
B 小块的列 -> N 方向

所以搬运 A/B 时，可能越两个方向。

A 搬运越界

int gr = blockIdx.y * BM + r;  // A 的全局行，M 方向
int gc = kk + c;               // A 的全局列，K 方向

sA[r][c] = A[gr * K + gc];

• gr >= M：M 方向最后一块伸出去了
• gc >= K：K 方向最后一片伸出去了

例如 K = 1000, BK = 16，最后一轮：

kk = 992
c = 0..15
gc = 992..1007

有效：992..999
越界：1000..1007

图示：

K = 1000，BK = 16，最后一轮 kk=992

A 的 K 方向：
[ ... 有效 992..999 ... ][ 1000..1007 越界 ]
                         └── 矩阵没有这些列

如果直接读 A[gr*K + gc]，gc>=K 会跳到下一行开头，读到别的数据；如果 gr 也是最后一行，就可能物理越界。

B 搬运越界

int gr = kk + r;               // B 的全局行，K 方向
int gc = blockIdx.x * BN + c;  // B 的全局列，N 方向

sB[r][c] = B[gr * N + gc];

• gr >= K：K 最后一片越界
• gc >= N：N 方向最后一块越界

正确做法：不读全局内存，但往 smem 写 0

sA[r][c] = (gr < M && gc < K) ? A[gr * K + gc] : 0.0f;
sB[r][c] = (gr < K && gc < N) ? B[gr * N + gc] : 0.0f;

注意这不是“跳过赋值”，而是：越界不读全局内存，但一定要往 smem 写 0。

为什么补 0 不影响数学结果？

矩阵乘法里，K 只到 K-1：

如果 K 不是 BK 整倍数，最后一轮多出来的 k 是：

k = K, K+1, ..., kk+BK-1

这些 k 本来就不该参与求和。把它们对应的 A/B 当作 0：

A[i][k] = 0  或  B[k][j] = 0

那么：

0 * B[k][j] = 0
A[i][k] * 0 = 0
0 * 0 = 0

加到 acc 上不改变结果：

acc += 0

图示：

K = 1000, BK = 16, 最后一轮 kk = 992

k:     992 ... 999 | 1000 ... 1007
A/B:   有效数据     | 补 0
求和:  正常加       | 加 0，不影响

为什么搬运越界“直接 continue 跳过不写 smem”是错的？

因为 shared memory 是每轮 K 循环复用的。

第 r 轮：
sA/sB 里是第 r 轮的数据

第 r+1 轮：
同样的 sA/sB 被覆盖

如果越界位置你直接跳过：

if (gr < M && gc < K) {
    sA[r][c] = A[gr * K + gc];
}
// 越界时什么都不写

那么 sA[r][c] 里留下的可能是：

第一轮：未初始化的垃圾
后面轮：上一轮 K 切片的旧数据

图示：

sA 的某个越界位置：

第 r 轮：写了有效值 5.0
第 r+1 轮：越界，跳过不写
          -> 仍然是 5.0

累加时：acc += 5.0 * sB[...]  // 垃圾参与计算

这就是“数据污染”：不崩不报错，但结果错。

正确做法是：

sA[r][c] = (gr < M && gc < K) ? A[gr * K + gc] : 0.0f;

无论有效还是越界，每个 smem 位置都必须被写一次：

有效 -> 写真实值
越界 -> 写 0

7.5 边界 3：累加环节数据污染

这不是一个新的“越界”，而是上一条没防住后的连锁反应。

如果搬运时越界读了垃圾，又没补 0，那么 sA/sB 里就有垃圾。累加时：

acc += sA[threadIdx.y][k] * sB[k][threadIdx.x];

垃圾就参与乘法，结果错。更阴险的是：它不一定崩，只是部分行列结果飘忽。

推荐第一版这样做：

搬运时越界 -> 写 0 进 smem
累加时固定 for k=0..BK-1
无效 k 位置乘到 0，不影响 acc

所以一般累加时不用再写 if (kk+k < K)。补 0 后统一累加，最简单、最不容易错。

如果你硬要跳过无效 k：

for (int k = 0; k < BK; ++k) {
    if (kk + k < K) {
        acc += ...;
    }
}

理论上也可以，但多了一层判断，而且没必要。

7.6 三处边界总表：谁 guard 谁，谁补 0，谁跳过

位置	代码	防什么	越界时怎么做
C 写回	if (row<M && col<N) C[...] = acc	写越界	直接跳过，不写
A 搬运	(gr<M && gc<K) ? A[...] : 0	读越界 + smem 垃圾	不读，但写 0 到 smem
B 搬运	(gr<K && gc<N) ? B[...] : 0	读越界 + smem 垃圾	不读，但写 0 到 smem
累加	固定 for k=0..BK-1	垃圾参与累加	通常不跳过，靠补 0 乘 0
__syncthreads()	无条件执行	死锁	绝对不能跳过

为什么搬运要补 0，而写回直接跳过？

搬运：smem 里每个位置都必须有值，否则累加会读到未初始化垃圾
      -> 越界不能读 A/B，但必须给 sA/sB 写一个 0

写回：C 的越界位置本来就不该存在
      -> 不是“写 0”，而是“根本不写”

图示：

搬运：
越界位置 -> sA/sB 写 0 -> 累加时乘 0，不影响 acc

写回：
越界位置 -> 什么都不做 -> C 不被污染

搬运 guard 和写回 guard 不是一回事，不能互相替代。

7.7 越界线程在整个 kernel 里到底干什么？

越界线程（比如 row >= M）不能提前 return，因为 __syncthreads() 要求全 block 所有线程都到达。它要这样走完全程：

搬运   -> 参与（越界位置往 smem 写 0）
屏障1  -> 到达
累加   -> 参与（算出一个没用的 acc）
屏障2  -> 到达
写回   -> 被 if 挡住，不写 C

图示：

越界线程：

搬运   -> 参与
屏障1  -> 到达
累加   -> 参与（acc 没用）
屏障2  -> 到达
写回   -> 被 if 挡住，不写

它只“白算”，但不“出错”。

为什么累加时越界线程不会读越界？

acc += sA[threadIdx.y][k] * sB[k][threadIdx.x];

它用的是 threadIdx.y 和 threadIdx.x 访问 shared memory：

sA[threadIdx.y][k]   -> 下标在 BM×BK 内，不越界
sB[k][threadIdx.x]   -> 下标在 BK×BN 内，不越界

所以越界线程只是算出一个无意义的 acc，但不会读越界、不会写越界。这个 acc 最后被 if 丢掉就行了。

7.8 常见错误写法对比

错误 1：C 写回不 guard

C[row * N + col] = acc;   // 越界线程写越界地址

后果：段错误，或悄悄改坏别的内存。

错误 2：把 guard 包到屏障外面

if (row < M && col < N) {
    __syncthreads();   // 死锁
}

后果：部分线程永远不到屏障，卡死。

错误 3：搬运越界直接 continue，不写 smem

if (gr < M && gc < K) {
    sA[r][c] = A[gr * K + gc];
}
// 越界时什么都不写

后果：sA/sB 残留上一轮旧值或未初始化垃圾，累加被污染。

错误 4：用写回的条件去 guard 搬运

// 错误理解：以为 row<M 就能保证搬运不越界
if (row < M) {
    sA[...] = A[...];
}

搬运的越界条件和写回不同：

写回看：row < M && col < N
搬运 A 看：gr < M && gc < K     （K 方向，不是 N）
搬运 B 看：gr < K && gc < N     （K 方向，不是 M）

而且搬运如果用 if 跳过而不补 0，sA/sB 会留下垃圾。

正确写法：

// 搬运：越界补 0，所有线程都参与
sA[r][c] = (gr < M && gc < K) ? A[gr * K + gc] : 0.0f;
sB[r][c] = (gr < K && gc < N) ? B[gr * N + gc] : 0.0f;

__syncthreads();   // 全员

// 累加：所有线程都算，越界线程的 acc 最后丢掉
for (int k = 0; k < BK; ++k) {
    acc += sA[threadIdx.y][k] * sB[k][threadIdx.x];
}

__syncthreads();   // 全员

// 写回：只有真实存在的格子才写
if (row < M && col < N) {
    C[row * N + col] = acc;
}

7.9 本节的速记句

越界 = 算出的全局坐标超出矩阵有效范围，却仍然去读写。

tile 固定 + grid 向上取整，最后一块天然会伸出边界。

C 写回越界直接跳过：不影响，因为那个格子本来就不该写。

A/B 搬运越界补 0：不影响，因为 0 乘任何有效数都是 0。

A/B 搬运越界直接跳过不写 smem：会影响，因为 smem 会残留上一轮或未初始化垃圾。

屏障不能跳过：条件里的 __syncthreads() 会死锁。

另外注意：原文件里写“K 非 BK 整倍数，如 K=1000, BK=8”，但 1000 / 8 = 125 是整除的，实际上不构成非整倍数。你实验时请选真正有余数的组合，比如 K=1001, BK=8，或 K=1000, BK=16。

8. shared memory 声明样板

起步用静态：

__shared__ float sA[BM][BK];
__shared__ float sB[BK][BN];

动态 smem 以后扫参再换：

extern __shared__ float buf[];
float* sA = buf;              // 前 BM*BK 个给 A
float* sB = buf + BM * BK;    // 后 BK*BN 个给 B

启动时第三个参数给字节数。今天第一版建议静态，中途换机制会导致扫描表前后数据不可比。

9. 动手任务怎么做

写 sgemm_tiled_v1.cu，接入台阶 01 harness。建议分两步：

第一步：只搬运 + 打印 smem

先让 kernel 只做“搬 A/B 进 smem + 打印”，对拍 Day 1 的分工图。确认每个线程搬的位置、索引、补零都对。

第二步：再加累加

搬运确认后再加 acc += sA * sB。一步一验，出错才知道在哪段找。

对拍形状：

形状	目的
1024³	常规整块
512×1024×256	非方阵，M/N/K 不同
1000×1000×1000	M/N 边界
K 非 BK 整倍数	K 边界，真正制造余数，如 K=1001, BK=8 或 K=1000, BK=16

PASS 后记录 GFLOPS、BM/BN/BK、两道屏障位置。只记录，不对比，不调优。

10. 出错排查顺序

卡住不动 / GPU 100% 无进展
  -> 搜 __syncthreads 是否在条件分支里

结果飘忽 / 部分行列错
  -> compute-sanitizer --tool racecheck   // 专抓 smem 竞态

段错误 / 写越界
  -> compute-sanitizer 默认 memcheck
  -> 查 C 写回 guard

结果全错或垃圾
  -> 先只搬运打印 smem，对 Day 1 分工图
  -> 查 A/B 索引算式、K 越界补零
  -> 特别查“搬运越界是不是直接 continue 没写 0”

11. 当日检验点，用口语回答

1. 四组形状全 PASS：1024³、512×1024×256、1000³、K 非 BK 整倍。
2. 能口头回答第二道屏障防什么：
   > 防第 r 轮算完后，快线程下一轮 load 覆盖 smem，而慢线程还在读第 r 轮 smem 的跨轮 WAR 竞态。
3. 第一版 GFLOPS 和两道屏障位置已留痕。
4. 没有把屏障放进任何条件分支，自查过。
5. 能口头说清楚三处边界：
   > C 写回越界直接跳过；A/B 搬运越界不读全局但必须写 0 到 smem；累加靠补 0 乘 0；屏障永远全员无条件。

最后一张速查图

每个 block：
  C_tile = 0

  for kk = 0; kk < K; kk += BK:
      所有线程搬 A 竖条 BM×BK -> sA
      所有线程搬 B 横条 BK×BN -> sB
      越界补 0（不是跳过不写）

      __syncthreads();     // 屏障1：搬完才准读

      每个线程：
          for k in 0..BK-1:
              acc += sA[ty][k] * sB[k][tx]

      __syncthreads();     // 屏障2：读完才准下一轮覆盖

  if (row < M && col < N):
      C[row*N+col] = acc     // 越界直接跳过，不写

记住这几句话：

搬运可以 if 挑人，屏障必须全员无条件。

屏障 1 防读垃圾，屏障 2 防下一轮覆盖。

C 写回要 guard，K 搬运要补零，读越界最阴险。

写回越界是“不写”，搬运越界是“写 0”；一个跳过，一个补零，别混。

如果后续还遇到 CUDA 底层的边角问题，也欢迎到 云栈社区 一起排查讨论。