第 8 课：朴素转置与合并访存

参考答案版：包含完整代码实现与问答题深度参考解析。建议先独立完成练习题后再展开核对。

本课目标与完成标准

学完后你应能：正确索引非方阵转置，并从 warp 地址序列解释跨步写。

通过必须同时满足：

独立补齐本课唯一的代码填空题，并通过给定检查；
三个问答题均说明因果链，而不是只报术语；
能指出至少一个正确性边界和一个性能取舍；
总分不低于 8/10，且没有一票否决级概念错误。

前置关系

课程前置：C/C++、线性代数、基本并行编程
本课在路线中的作用：矩阵转置交换行列；row-major 输入连续读后，直接写转置位置通常形成大步长 store。

核心心智模型

1. 它是什么，解决什么问题

矩阵转置（Matrix Transpose） 是高性能线性代数、自注意力多头变换（如 $B \times S \times H \times D \to B \times H \times S \times D$ ）以及卷积 Im2Col 中最常见的数据重排操作。数学上，将行优先存储的矩阵 $A \in \mathbb{R}^{M \times N}$ 转换为矩阵 $B = A^T \in \mathbb{R}^{N \times M}$ ： $B[j, i] = A[i, j], \quad 0 \le i < M, \ 0 \le j < N$

底层硬件挑战（访存不对称性与总线击穿）： - 在行优先（Row-Major）布局中， $A[i, j]$ 在内存中的物理线性偏移量为 $\text{offset}_A = i \cdot N + j$ ，而转置矩阵 $B[j, i]$ 的物理线性偏移量为 $\text{offset}_B = j \cdot M + i$ （注意输出的主维度步长 Leading Dimension 变成了 $M$ ）； - 当 GPU 线程并发执行时，即使读端做到了相邻线程访问连续内存（完全合并访存），写端也会因为地址跨步为 $M$ 而将连续的 Warp 写操作打散到 32 个不同的内存扇区中，导致显存总线事务数暴增 8 倍，有效带宽急剧下跌（图 8-1 与图 8-2）。

2. 它如何工作（结合图 8-1 与图 8-2）

朴素矩阵转置（Naive Transpose）采用 One Thread per Element（一个线程负责一个元素的搬运）：

全局坐标映射：
- 启动 2D Grid 和 2D Block（如 dim3 block(16, 16)）；
- 线程全局坐标分别为列索引 col = blockIdx.x * blockDim.x + threadIdx.x 与行索引 row = blockIdx.y * blockDim.y + threadIdx.y；
- 执行边界守卫：当且仅当 $row < M$ 且 $col < N$ 时线程参与计算。
读操作：完美合并读取（Coalesced Load）：
- 同一个 Warp 内的 32 个线程具有相同的 row，但 col 连续递增（ $col, col+1, \dots, col+31$ ）；
- 读取地址为 $row \cdot N + col$ ，32 个线程请求的物理地址在显存中完全连续，总计 $32 \times 4\text{ B} = 128\text{ B}$ ；
- 硬件显存控制器仅需发起 1 次 128 字节 DRAM 突发事务，显存总线利用率达到 100%。
写操作：大步长非合并写入（Strided Store，图 8-1 与图 8-2）：
- 目标地址为 $col \cdot M + row$ 。由于相邻线程的 col 连续递增，写入地址在内存中的跨度为 $M$ 个浮点数（即 $4M$ 字节）；
- 当 $M \ge 32$ 时，32 个线程的写入地址分别落入 32 个完全不同的 32 字节扇区或 128 字节 Cache Line；
- 显存控制器被迫发起 32 次独立的内存写事务。每次传输 32 字节却只写入 4 字节有效数据，有效显存写入带宽直接萎缩到理论峰值的 $\frac{4}{32} = 12.5\%$ ！

flowchart TD
    subgraph CoalescedRead["读端: 完美合并读取 (Coalesced Load, 100% 效率)"]
        direction LR
        Warp_R["Warp 32 连续线程 (Lane 0..31)"] -->|"请求连续地址: row*N + col..col+31"| Burst["单次 128B DRAM 突发事务 (100% 满载)"]
    end

    subgraph StridedWrite["写端: 严重跨步非合并写 (Strided Write, 12.5% 效率)"]
        direction LR
        Warp_W["Warp 32 连续线程 (Lane 0..31)"] -->|"跨度步长为 M (每次跳跃 4M 字节)"| Split["打散为 32 次独立的 32B 扇区事务 (严重总线停顿)"]
    end

    CoalescedRead -.->|转置数据流向| StridedWrite

图 8-1：朴素矩阵转置读端合并访存 vs 写端跨步离散访存微架构总线对比

flowchart LR
    In["输入矩阵 A[row, col]<br>Shape: M × N, Leading Dim = N"] --> Thread["单线程 (row, col) 读取并对调索引"]
    Thread --> Out["输出矩阵 B[col, row]<br>Shape: N × M, Leading Dim = M"]

图 8-2：朴素矩阵转置单个线程对调逻辑坐标与 Leading Dimension 变更流程

3. 正确性条件与常见误区

行长与列长的严格区分（ $M$ 与 $N$ 的维度混淆）：
- 输入矩阵的 Leading Dimension 是 $N$ （一行有 $N$ 个元素），物理索引为 $row \cdot N + col$ ；
- 输出矩阵的 Leading Dimension 变为 $M$ （一行有 $M$ 个元素），物理索引必须严格为 $col \cdot M + row$ ，而不是 $col \cdot N + row$ ；
- 边界判断必须严格为 row < M && col < N，不可写反。
测试用例设计的方阵陷阱：
- 如果仅在 $M = N$ 的方阵（如 $1024 \times 1024$ ）上测试，上述两处致命的索引倒置与边界混淆在数值上完全等价（ $M=N$ ），Bug 会被 100% 掩盖。必须使用非对称长方阵（如 $1024 \times 512$ ）才能暴露索引缺陷。

4. 性能与工程取舍

朴素搬运 vs. Shared Memory Tiling 缓冲：
- 朴素版：Kernel 代码极其简洁，无任何线程间通信与同步，零 Shared Memory 开销；但在长行长列场景下，写带宽利用率跌入谷底，严重受制于显存控制器并发上限；
- Tiled 共享内存版（下一课）：在片上分配二维 Shared Memory 瓦片。读入时保持合并读写入 SMEM；通过在 SMEM 内部进行索引对调（转置），再从 SMEM 读出合并写回 Global Memory。用少量片上同步换取读写两端的 100% 合并，性能提升数倍。

具体演示：非方阵物理地址步长对比实测

考虑矩阵维度 $M = 2$ （行数）， $N = 3$ （列数），元素以浮点数存储（每个 4 字节）：

输入矩阵 $A_{2 \times 3}$ ： $\begin{bmatrix} A_{0,0} & A_{0,1} & A_{0,2} \\ A_{1,0} & A_{1,1} & A_{1,2} \end{bmatrix}$ 物理存储线性地址： $[0, 1, 2, 3, 4, 5]$ ，步长 $\Delta \text{col} = 1$ 个元素。
输出转置矩阵 $B_{3 \times 2}$ ： $\begin{bmatrix} B_{0,0} & B_{0,1} \\ B_{1,0} & B_{1,1} \\ B_{2,0} & B_{2,1} \end{bmatrix}$ 输出矩阵每行只有 $M = 2$ 个元素。
- 线程 $(row=0, col=0)$ 搬运 $A_{0,0} \to B_{0,0}$ ，输出线性索引 $0 \cdot 2 + 0 = 0$ ；
- 同 Warp 相邻线程 $(row=0, col=1)$ 搬运 $A_{0,1} \to B_{1,0}$ ，输出线性索引 $1 \cdot 2 + 0 = 2$ ；
- 相邻线程 $(row=0, col=2)$ 搬运 $A_{0,2} \to B_{2,0}$ ，输出线性索引 $2 \cdot 2 + 0 = 4$ ；
微观因果：相邻线程在写端产生的地址间距为 $M = 2$ 个元素（8 字节）。随着 $M$ 扩大到 1024 或 4096，相邻线程写的地址跨度达到数 KB，完全粉碎了 DRAM 的连续突发特性。

实践任务：唯一代码填空题

补齐输出线性索引。

规则：只能修改 TODO/______ 所在位置；不要删除断言或放宽误差。代码注释说明了每个边界条件。

Code

%%writefile /tmp/08_transpose_naive.cu
#include <cuda_runtime.h>

// Naive Transpose:
// input:  [M, N] row-major
// output: [N, M] row-major
//
// 这个版本一个线程搬一个元素。
// 读 input 时，相邻线程读 input[row, col] 的连续地址，读是合并的；
// 写 output[col, row] 时，相邻线程写的地址跨度是 M，写通常不合并。
//
// 面试讲法：
// naive transpose 的瓶颈通常在 strided global store，
// 优化版使用 shared memory tile，把 strided store 变成 coalesced store。
__global__ void transpose_naive_kernel(
    const float* __restrict__ input,
    float* __restrict__ output,
    int M,
    int N
) {
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;

    if (row < M && col < N) {
        output[______] = input[row * N + col];  // TODO: 非方阵也正确
    }
}

void launch_transpose_naive(
    const float* input,
    float* output,
    int M,
    int N,
    cudaStream_t stream
) {
    dim3 block(16, 16);
    dim3 grid((N + block.x - 1) / block.x, (M + block.y - 1) / block.y);
    transpose_naive_kernel<<<grid, block, 0, stream>>>(input, output, M, N);
}

检查方法

有 CUDA 环境时执行 nvcc -std=c++17 -c /tmp/08_transpose_naive.cu -o /tmp/08_transpose_naive.cu.o；无 CUDA 环境时只做静态审查并登记待验证。

提交时请给出：补齐后的代码、实际运行输出（环境不可用时注明“仅静态审查”）以及对失败用例的解释。

Q1

不要背定义：请从输入、状态变化和输出三个阶段解释“朴素转置与合并访存”的工作机制。

你的答案：

Q2

为何只用 1024×1024 方阵测试容易漏掉索引 bug？

你的答案：

Q3

列主序输入时这套索引和合并访存结论如何变化？

你的答案：

评分与通过规则

代码 4 分：正常输入 2 分，边界输入 1 分，解释实现 1 分；
Q1～Q3 各 2 分；
一票否决：结果碰巧正确但核心因果链错误、删除边界检查、把未运行结果说成实测。

需要提示时按四级机制请求：概念区域 → 具体方向 → 关键局部 → 完整答案。

💡 点击展开参考答案与详细代码实现

参考答案（仅 answer 分支）

先完成题目再核对。即使代码一致，也要能解释关键步骤，并尝试更换一个输入规模。

Code

%%writefile /tmp/08_transpose_naive.cu
#include <cuda_runtime.h>

// Naive Transpose:
// input:  [M, N] row-major
// output: [N, M] row-major
//
// 这个版本一个线程搬一个元素。
// 读 input 时，相邻线程读 input[row, col] 的连续地址，读是合并的；
// 写 output[col, row] 时，相邻线程写的地址跨度是 M，写通常不合并。
//
// 面试讲法：
// naive transpose 的瓶颈通常在 strided global store，
// 优化版使用 shared memory tile，把 strided store 变成 coalesced store。
__global__ void transpose_naive_kernel(
    const float* __restrict__ input,
    float* __restrict__ output,
    int M,
    int N
) {
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;

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

void launch_transpose_naive(
    const float* input,
    float* output,
    int M,
    int N,
    cudaStream_t stream
) {
    dim3 block(16, 16);
    dim3 grid((N + block.x - 1) / block.x, (M + block.y - 1) / block.y);
    transpose_naive_kernel<<<grid, block, 0, stream>>>(input, output, M, N);
}

Q1 参考答案

输入阶段（Input Coordinate Mapping & Coalesced Load）：
- 启动二维网格与二维线程块（dim3 block(16, 16)），每个线程通过其二维 ID 计算在输入矩阵中的坐标：列 col = blockIdx.x * 16 + threadIdx.x，行 row = blockIdx.y * 16 + threadIdx.y；
- 同一个 Warp 内连续的 32 个线程拥有相同的 row，其 col 连续递增；线程通过行主序公式 row * N + col 计算读取地址。32 个线程请求的 128 字节物理显存地址完全连续，由内存控制器合并为 1 次 128 字节 DRAM 突发传输完成加载；
状态变化阶段（Register Staging & Coordinate Swap）：
- 数据被短暂暂存在线程本地私有寄存器中；
- 线程在逻辑上将自身的矩阵坐标对调：原输入坐标 $(row, col)$ 对应转置矩阵的目标坐标 $(col, row)$ ；
输出阶段（Strided Store & Bus Fragmentation）：
- 输出矩阵由于行数列数对调，其主维度步长（Leading Dimension）变成了原输入行数 $M$ ；
- 目标物理写入地址为 col * M + row。由于同一个 Warp 中相邻线程的 col 连续递增，写入地址之间的物理跨度高达 $M$ 个浮点数（ $4M$ 字节）；
- 当 $M \ge 32$ 时，相邻线程写入落入完全不相交的内存扇区，导致写请求彻底无法合并，退化为 32 次独立的散落写事务，造成显存总线严重拥堵。

Q2 参考答案

仅使用 $1024 \times 1024$ 等对称方阵进行验证极易掩盖两类最致命的索引与边界 Bug：

Leading Dimension 步长混淆被完全屏蔽：
- 正确的输出索引公式是 col * M + row（因为输出矩阵每行长度为 $M$ ）；
- 如果程序员思维惯性误写成 col * N + row，或者将输入输出步长变量颠倒：
- 当 $M = N = 1024$ 时， $M$ 与 $N$ 数值完全相同，col * M + row 与 col * N + row 计算出的线性偏移量 100% 一模一样！测试用例的误差比对不仅不会报错，甚至能拿到满分 pass；
长宽边界守卫混淆被完全屏蔽：
- 若程序员将越界保护误写为 row < N && col < M（把行上限错用为列数，列上限错用为行数）；
- 在方阵测试下， $M = N$ ，该边界条件依然成立，显存越界访问和线程退出逻辑无法被任何断言检测到；
- 一旦上线部署到非方阵业务（例如 Prefill 阶段常见矩阵尺寸如 $M=4096, N=128$ ），错误的索引公式会引发致命的显存越界段错误（cudaErrorIllegalAddress）或严重的权重数据错位污染。因此单元测试必须覆盖非方阵场景。

Q3 参考答案

若输入矩阵为列主序（Column-Major，例如 Fortran 或某些 BLAS 库的数据布局），读写两端的合并访存结论将发生完全的对偶反转：

读写索引公式重构：
- 输入矩阵为列主序时，元素 $A[row, col]$ 的物理内存地址为 col * M + row（一列有 $M$ 个连续元素）；
- 输出矩阵若转置为行主序，目标地址为 col * M + row；若输出依然保持列主序，目标地址为 row * N + col；
合并访存瓶颈由写端转移到读端：
- 同一个 Warp 内连续的 32 个线程（拥有连续的 threadIdx.x，对应连续的 col）在读取输入时，地址变为 col * M + row；
- 此时读操作变成了大步长跨步访存（Strided Load），相邻线程读取的物理地址跨度为 $M$ 个元素，读请求无法合并，每次读入有效载荷只有 12.5%；
- 相反，如果写端写入连续行主序地址，写操作反而变成了完全合并写入；
- 核心启示：无论数据是行主序还是列主序，朴素的转置搬运必然导致“一端合并、另一端跨步”的硬件物理矛盾。要消除这一矛盾，必须借助片上 Shared Memory 进行二维数据重排（下一课）。

参考资料

资料用于建立事实基线；面试回答仍需用自己的语言组织。

← 上一课
第 7 课：Online Softmax 与… 下一课 →
第 9 课：共享内存转置与 Bank Con…

--- title: "第 8 课：朴素转置与合并访存" jupyter: python3 --- ::: {.callout-note appearance="simple"} **参考答案版**：包含完整代码实现与问答题深度参考解析。建议先独立完成练习题后再展开核对。 ::: ## 本课目标与完成标准学完后你应能：正确索引非方阵转置，并从 warp 地址序列解释跨步写。通过必须同时满足： - 独立补齐本课唯一的代码填空题，并通过给定检查； - 三个问答题均说明因果链，而不是只报术语； - 能指出至少一个正确性边界和一个性能取舍； - 总分不低于 8/10，且没有一票否决级概念错误。 ## 前置关系 - 课程前置：C/C++、线性代数、基本并行编程 - 本课在路线中的作用：矩阵转置交换行列；row-major 输入连续读后，直接写转置位置通常形成大步长 store。 ## 核心心智模型 ### 1. 它是什么，解决什么问题 **矩阵转置（Matrix Transpose）** 是高性能线性代数、自注意力多头变换（如 $B \times S \times H \times D \to B \times H \times S \times D$）以及卷积 Im2Col 中最常见的数据重排操作。数学上，将行优先存储的矩阵 $A \in \mathbb{R}^{M \times N}$ 转换为矩阵 $B = A^T \in \mathbb{R}^{N \times M}$： $$B[j, i] = A[i, j], \quad 0 \le i < M, \ 0 \le j < N$$ **底层硬件挑战（访存不对称性与总线击穿）**： - 在行优先（Row-Major）布局中，$A[i, j]$ 在内存中的物理线性偏移量为 $\text{offset}_A = i \cdot N + j$，而转置矩阵 $B[j, i]$ 的物理线性偏移量为 $\text{offset}_B = j \cdot M + i$（注意输出的主维度步长 Leading Dimension 变成了 $M$）； - 当 GPU 线程并发执行时，即使读端做到了相邻线程访问连续内存（完全合并访存），写端也会因为地址跨步为 $M$ 而将连续的 Warp 写操作打散到 32 个不同的内存扇区中，导致显存总线事务数暴增 8 倍，有效带宽急剧下跌（图 8-1 与图 8-2）。 ### 2. 它如何工作（结合图 8-1 与图 8-2）朴素矩阵转置（Naive Transpose）采用 **One Thread per Element（一个线程负责一个元素的搬运）**： 1. **全局坐标映射**： - 启动 2D Grid 和 2D Block（如 `dim3 block(16, 16)`）； - 线程全局坐标分别为列索引 `col = blockIdx.x * blockDim.x + threadIdx.x` 与行索引 `row = blockIdx.y * blockDim.y + threadIdx.y`； - 执行边界守卫：当且仅当 $row < M$ 且 $col < N$ 时线程参与计算。 2. **读操作：完美合并读取（Coalesced Load）**： - 同一个 Warp 内的 32 个线程具有相同的 `row`，但 `col` 连续递增（$col, col+1, \dots, col+31$）； - 读取地址为 $row \cdot N + col$，32 个线程请求的物理地址在显存中完全连续，总计 $32 \times 4\text{ B} = 128\text{ B}$； - 硬件显存控制器仅需发起 **1 次 128 字节 DRAM 突发事务**，显存总线利用率达到 100%。 3. **写操作：大步长非合并写入（Strided Store，图 8-1 与图 8-2）**： - 目标地址为 $col \cdot M + row$。由于相邻线程的 `col` 连续递增，写入地址在内存中的跨度为 $M$ 个浮点数（即 $4M$ 字节）； - 当 $M \ge 32$ 时，32 个线程的写入地址分别落入 32 个完全不同的 32 字节扇区或 128 字节 Cache Line； - 显存控制器被迫发起 **32 次独立的内存写事务**。每次传输 32 字节却只写入 4 字节有效数据，有效显存写入带宽直接萎缩到理论峰值的 $\frac{4}{32} = 12.5\%$！ ```{mermaid} flowchart TD subgraph CoalescedRead["读端: 完美合并读取 (Coalesced Load, 100% 效率)"] direction LR Warp_R["Warp 32 连续线程 (Lane 0..31)"] -->|"请求连续地址: row*N + col..col+31"| Burst["单次 128B DRAM 突发事务 (100% 满载)"] end subgraph StridedWrite["写端: 严重跨步非合并写 (Strided Write, 12.5% 效率)"] direction LR Warp_W["Warp 32 连续线程 (Lane 0..31)"] -->|"跨度步长为 M (每次跳跃 4M 字节)"| Split["打散为 32 次独立的 32B 扇区事务 (严重总线停顿)"] end CoalescedRead -.->|转置数据流向| StridedWrite ``` 图 8-1：朴素矩阵转置读端合并访存 vs 写端跨步离散访存微架构总线对比 ```{mermaid} flowchart LR In["输入矩阵 A[row, col] Shape: M × N, Leading Dim = N"] --> Thread["单线程 (row, col) 读取并对调索引"] Thread --> Out["输出矩阵 B[col, row] Shape: N × M, Leading Dim = M"] ``` 图 8-2：朴素矩阵转置单个线程对调逻辑坐标与 Leading Dimension 变更流程 ### 3. 正确性条件与常见误区 1. **行长与列长的严格区分（$M$ 与 $N$ 的维度混淆）**： - 输入矩阵的 Leading Dimension 是 $N$（一行有 $N$ 个元素），物理索引为 $row \cdot N + col$； - 输出矩阵的 Leading Dimension 变为 $M$（一行有 $M$ 个元素），物理索引必须严格为 $col \cdot M + row$，而不是 $col \cdot N + row$； - 边界判断必须严格为 `row < M && col < N`，不可写反。 2. **测试用例设计的方阵陷阱**： - 如果仅在 $M = N$ 的方阵（如 $1024 \times 1024$）上测试，上述两处致命的索引倒置与边界混淆在数值上完全等价（$M=N$），Bug 会被 100% 掩盖。必须使用非对称长方阵（如 $1024 \times 512$）才能暴露索引缺陷。 ### 4. 性能与工程取舍 1. **朴素搬运 vs. Shared Memory Tiling 缓冲**： - **朴素版**：Kernel 代码极其简洁，无任何线程间通信与同步，零 Shared Memory 开销；但在长行长列场景下，写带宽利用率跌入谷底，严重受制于显存控制器并发上限； - **Tiled 共享内存版（下一课）**：在片上分配二维 Shared Memory 瓦片。读入时保持合并读写入 SMEM；通过在 SMEM 内部进行索引对调（转置），再从 SMEM 读出合并写回 Global Memory。用少量片上同步换取读写两端的 100% 合并，性能提升数倍。 ## 具体演示：非方阵物理地址步长对比实测考虑矩阵维度 $M = 2$（行数），$N = 3$（列数），元素以浮点数存储（每个 4 字节）： - **输入矩阵 $A_{2 \times 3}$**： $$\begin{bmatrix} A_{0,0} & A_{0,1} & A_{0,2} \\ A_{1,0} & A_{1,1} & A_{1,2} \end{bmatrix}$$ 物理存储线性地址：$[0, 1, 2, 3, 4, 5]$，步长 $\Delta \text{col} = 1$ 个元素。 - **输出转置矩阵 $B_{3 \times 2}$**： $$\begin{bmatrix} B_{0,0} & B_{0,1} \\ B_{1,0} & B_{1,1} \\ B_{2,0} & B_{2,1} \end{bmatrix}$$ 输出矩阵每行只有 $M = 2$ 个元素。 - 线程 $(row=0, col=0)$ 搬运 $A_{0,0} \to B_{0,0}$，输出线性索引 $0 \cdot 2 + 0 = 0$； - 同 Warp 相邻线程 $(row=0, col=1)$ 搬运 $A_{0,1} \to B_{1,0}$，输出线性索引 $1 \cdot 2 + 0 = 2$； - 相邻线程 $(row=0, col=2)$ 搬运 $A_{0,2} \to B_{2,0}$，输出线性索引 $2 \cdot 2 + 0 = 4$； - **微观因果**：相邻线程在写端产生的地址间距为 $M = 2$ 个元素（8 字节）。随着 $M$ 扩大到 1024 或 4096，相邻线程写的地址跨度达到数 KB，完全粉碎了 DRAM 的连续突发特性。 ## 实践任务：唯一代码填空题补齐输出线性索引。规则：只能修改 `TODO`/`______` 所在位置；不要删除断言或放宽误差。代码注释说明了每个边界条件。 ```{python} #| course_role: exercise %%writefile /tmp/08_transpose_naive.cu #include <cuda_runtime.h> // Naive Transpose: // input: [M, N] row-major // output: [N, M] row-major // // 这个版本一个线程搬一个元素。 // 读 input 时，相邻线程读 input[row, col] 的连续地址，读是合并的； // 写 output[col, row] 时，相邻线程写的地址跨度是 M，写通常不合并。 // // 面试讲法： // naive transpose 的瓶颈通常在 strided global store， // 优化版使用 shared memory tile，把 strided store 变成 coalesced store。 __global__ void transpose_naive_kernel( const float* __restrict__ input, float* __restrict__ output, int M, int N ) { int col = blockIdx.x * blockDim.x + threadIdx.x; int row = blockIdx.y * blockDim.y + threadIdx.y; if (row < M && col < N) { output[______] = input[row * N + col]; // TODO: 非方阵也正确 } } void launch_transpose_naive( const float* input, float* output, int M, int N, cudaStream_t stream ) { dim3 block(16, 16); dim3 grid((N + block.x - 1) / block.x, (M + block.y - 1) / block.y); transpose_naive_kernel<<<grid, block, 0, stream>>>(input, output, M, N); } ``` ### 检查方法有 CUDA 环境时执行 `nvcc -std=c++17 -c /tmp/08_transpose_naive.cu -o /tmp/08_transpose_naive.cu.o`；无 CUDA 环境时只做静态审查并登记待验证。提交时请给出：补齐后的代码、实际运行输出（环境不可用时注明“仅静态审查”）以及对失败用例的解释。 ### Q1 不要背定义：请从输入、状态变化和输出三个阶段解释“朴素转置与合并访存”的工作机制。 **你的答案：** ### Q2 为何只用 1024×1024 方阵测试容易漏掉索引 bug？ **你的答案：** ### Q3 列主序输入时这套索引和合并访存结论如何变化？ **你的答案：** ## 评分与通过规则 - 代码 4 分：正常输入 2 分，边界输入 1 分，解释实现 1 分； - Q1～Q3 各 2 分； - 一票否决：结果碰巧正确但核心因果链错误、删除边界检查、把未运行结果说成实测。需要提示时按四级机制请求：概念区域 → 具体方向 → 关键局部 → 完整答案。 ::: {.callout-tip collapse="true" title="💡 点击展开参考答案与详细代码实现" icon="false"} ## 参考答案（仅 answer 分支）先完成题目再核对。即使代码一致，也要能解释关键步骤，并尝试更换一个输入规模。 ```{python} #| course_role: solution %%writefile /tmp/08_transpose_naive.cu #include <cuda_runtime.h> // Naive Transpose: // input: [M, N] row-major // output: [N, M] row-major // // 这个版本一个线程搬一个元素。 // 读 input 时，相邻线程读 input[row, col] 的连续地址，读是合并的； // 写 output[col, row] 时，相邻线程写的地址跨度是 M，写通常不合并。 // // 面试讲法： // naive transpose 的瓶颈通常在 strided global store， // 优化版使用 shared memory tile，把 strided store 变成 coalesced store。 __global__ void transpose_naive_kernel( const float* __restrict__ input, float* __restrict__ output, int M, int N ) { int col = blockIdx.x * blockDim.x + threadIdx.x; int row = blockIdx.y * blockDim.y + threadIdx.y; if (row < M && col < N) { output[col * M + row] = input[row * N + col]; } } void launch_transpose_naive( const float* input, float* output, int M, int N, cudaStream_t stream ) { dim3 block(16, 16); dim3 grid((N + block.x - 1) / block.x, (M + block.y - 1) / block.y); transpose_naive_kernel<<<grid, block, 0, stream>>>(input, output, M, N); } ``` ### Q1 参考答案 1. **输入阶段（Input Coordinate Mapping & Coalesced Load）**： - 启动二维网格与二维线程块（`dim3 block(16, 16)`），每个线程通过其二维 ID 计算在输入矩阵中的坐标：列 `col = blockIdx.x * 16 + threadIdx.x`，行 `row = blockIdx.y * 16 + threadIdx.y`； - 同一个 Warp 内连续的 32 个线程拥有相同的 `row`，其 `col` 连续递增；线程通过行主序公式 `row * N + col` 计算读取地址。32 个线程请求的 128 字节物理显存地址完全连续，由内存控制器合并为 1 次 128 字节 DRAM 突发传输完成加载； 2. **状态变化阶段（Register Staging & Coordinate Swap）**： - 数据被短暂暂存在线程本地私有寄存器中； - 线程在逻辑上将自身的矩阵坐标对调：原输入坐标 $(row, col)$ 对应转置矩阵的目标坐标 $(col, row)$； 3. **输出阶段（Strided Store & Bus Fragmentation）**： - 输出矩阵由于行数列数对调，其主维度步长（Leading Dimension）变成了原输入行数 $M$； - 目标物理写入地址为 `col * M + row`。由于同一个 Warp 中相邻线程的 `col` 连续递增，写入地址之间的物理跨度高达 $M$ 个浮点数（$4M$ 字节）； - 当 $M \ge 32$ 时，相邻线程写入落入完全不相交的内存扇区，导致写请求彻底无法合并，退化为 32 次独立的散落写事务，造成显存总线严重拥堵。 ### Q2 参考答案仅使用 $1024 \times 1024$ 等对称方阵进行验证极易掩盖两类最致命的索引与边界 Bug： 1. **Leading Dimension 步长混淆被完全屏蔽**： - 正确的输出索引公式是 `col * M + row`（因为输出矩阵每行长度为 $M$）； - 如果程序员思维惯性误写成 `col * N + row`，或者将输入输出步长变量颠倒： - 当 $M = N = 1024$ 时，$M$ 与 $N$ 数值完全相同，`col * M + row` 与 `col * N + row` 计算出的线性偏移量 100% 一模一样！测试用例的误差比对不仅不会报错，甚至能拿到满分 pass； 2. **长宽边界守卫混淆被完全屏蔽**： - 若程序员将越界保护误写为 `row < N && col < M`（把行上限错用为列数，列上限错用为行数）； - 在方阵测试下，$M = N$，该边界条件依然成立，显存越界访问和线程退出逻辑无法被任何断言检测到； - 一旦上线部署到非方阵业务（例如 Prefill 阶段常见矩阵尺寸如 $M=4096, N=128$），错误的索引公式会引发致命的显存越界段错误（`cudaErrorIllegalAddress`）或严重的权重数据错位污染。因此单元测试必须覆盖非方阵场景。 ### Q3 参考答案若输入矩阵为列主序（Column-Major，例如 Fortran 或某些 BLAS 库的数据布局），读写两端的合并访存结论将发生**完全的对偶反转**： 1. **读写索引公式重构**： - 输入矩阵为列主序时，元素 $A[row, col]$ 的物理内存地址为 `col * M + row`（一列有 $M$ 个连续元素）； - 输出矩阵若转置为行主序，目标地址为 `col * M + row`；若输出依然保持列主序，目标地址为 `row * N + col`； 2. **合并访存瓶颈由写端转移到读端**： - 同一个 Warp 内连续的 32 个线程（拥有连续的 `threadIdx.x`，对应连续的 `col`）在读取输入时，地址变为 `col * M + row`； - 此时**读操作变成了大步长跨步访存（Strided Load）**，相邻线程读取的物理地址跨度为 $M$ 个元素，读请求无法合并，每次读入有效载荷只有 12.5%； - 相反，如果写端写入连续行主序地址，写操作反而变成了完全合并写入； - **核心启示**：无论数据是行主序还是列主序，朴素的转置搬运必然导致“一端合并、另一端跨步”的硬件物理矛盾。要消除这一矛盾，必须借助片上 Shared Memory 进行二维数据重排（下一课）。 ## 参考资料 - [CUDA Programming Guide](https://docs.nvidia.com/cuda/cuda-programming-guide/) - [CUDA Best Practices Guide](https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/) 资料用于建立事实基线；面试回答仍需用自己的语言组织。 ::: ::: {.course-pagination} [**← 上一课** 第 7 课：Online Softmax 与...](lesson07_online_softmax.qmd){.pagination-link .pagination-prev} [**下一课 →** 第 9 课：共享内存转置与 Bank Con...](lesson09_transpose_tiled.qmd){.pagination-link .pagination-next} :::