Naive GEMM 的致命问题：每个 thread 计算 C 的一个元素时， 要从 Global Memory 读 K 次 A 和 K 次 B—— 而相邻 thread 读的数据大量重叠。 Tiling 的核心思想：把数据分块搬进 Shared Memory，让整个 Block 的 thread 复用同一份数据。

Tiling 解决了数据复用问题，但还有两个关键的访存陷阱： Global Memory 的 coalesced access 和 Shared Memory 的 bank conflict。 不注意这两点，tiling 的加速效果会大打折扣。

Nsight Compute（ncu）是 NVIDIA 的 kernel-level profiler， 能精确量化 occupancy、memory throughput、compute throughput， 告诉你 kernel 是 compute-bound 还是 memory-bound。

常用命令

# 基础 profiling — 收集所有指标
ncu --set full -o report ./tiled_gemm

# 只看 memory 相关指标
ncu --metrics l1tex__t_bytes_pipe_lsu_mem_global_op_ld.sum,\
l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum,\
smsp__sass_average_data_bytes_per_sector_mem_global_op_ld.pct \
./tiled_gemm

# 对比两个 kernel
ncu --set full -o naive  ./naive_gemm
ncu --set full -o tiled  ./tiled_gemm
# 用 Nsight Compute GUI 打开 .ncu-rep 文件对比

# 看 occupancy 限制因素
ncu --metrics sm__warps_active.avg.pct_of_peak_sustained_active \
./tiled_gemm

Lab — tiled_gemm.cu

在 Day 04 的 naive GEMM 基础上，加入 Shared Memory tiling。 每个 Block 负责 C 的一个 TILE_SIZE×TILE_SIZE 子矩阵， 通过循环加载 A 和 B 的 tile 到 SMEM，计算部分积，最终累加写回。

// tiled_gemm.cu — Shared Memory Tiled Matrix Multiply
#define TILE_SIZE 16

__global__ void tiled_gemm(float *A, float *B, float *C,
                            int M, int N, int K) {
    // Block 内的线程坐标
    int tx = threadIdx.x, ty = threadIdx.y;
    // 全局输出坐标
    int row = blockIdx.y * TILE_SIZE + ty;
    int col = blockIdx.x * TILE_SIZE + tx;

    // Shared Memory (padding 消除 bank conflict)
    __shared__ float As[TILE_SIZE][TILE_SIZE + 1];
    __shared__ float Bs[TILE_SIZE][TILE_SIZE + 1];

    float sum = 0.0f;

    // 沿 K 维度分块循环
    for (int t = 0; t < (K + TILE_SIZE - 1) / TILE_SIZE; t++) {
        // 1. 协作加载 tile 到 SMEM (coalesced!)
        int aCol = t * TILE_SIZE + tx;
        int bRow = t * TILE_SIZE + ty;
        As[ty][tx] = (row < M && aCol < K) ? A[row * K + aCol] : 0.0f;
        Bs[ty][tx] = (bRow < K && col < N) ? B[bRow * N + col] : 0.0f;

        // 2. 同步：确保 tile 加载完毕
        __syncthreads();

        // 3. 从 SMEM 计算部分积
        for (int k = 0; k < TILE_SIZE; k++) {
            sum += As[ty][k] * Bs[k][tx];
        }

        // 4. 同步：确保计算完毕再加载下一个 tile
        __syncthreads();
    }

    // 写回结果
    if (row < M && col < N) {
        C[row * N + col] = sum;
    }
}

// 启动配置
dim3 block(TILE_SIZE, TILE_SIZE);     // 16×16 = 256 threads
dim3 grid((N + TILE_SIZE - 1) / TILE_SIZE,
          (M + TILE_SIZE - 1) / TILE_SIZE);
tiled_gemm<<<grid, block>>>(d_A, d_B, d_C, M, N, K);

性能对比实验

用 1024×1024 矩阵，分别跑 naive GEMM、tiled GEMM、cuBLAS，记录 GFLOPS。

1. 画出 Tiled GEMM 的数据流：Global Memory → Shared Memory → Register → 写回 Global Memory，标注 __syncthreads() 的位置。
2. Naive GEMM (1024×1024) 的 Global Memory 读取量是多少字节？Tiled GEMM (T=16) 呢？计算降低比例。
3. 解释为什么 __shared__ float As[16][16] 列方向读取有 bank conflict，而 As[16][17] 没有。
4. 你的 tiled GEMM 在 ncu 中的 Achieved Occupancy 是多少？限制因素是什么（registers / shared memory / block size）？
5. 对比你的 naive、tiled、cuBLAS 三个版本的 GFLOPS，tiling 带来了多少倍加速？距离 cuBLAS 还差多少？