概述Tiling分块把一个大计算或大张量切成许多小块每块单独计算让数据能装进片上高速存储。例子餐厅比喻Naive不分块 从仓库HBM拿一袋米 → 跑到厨房SM→ 用一勺 → 跑回仓库放好 → 再拿一袋 → 再跑 → 再用一勺 → ... 每次用一勺米都要跑一趟仓库累死。 Tiled分块 从仓库一次搬一整袋米Tile到厨房台面Shared Memory → 在台面上反复用Compute → 用完了再搬下一袋 跑仓库的次数大幅减少。具体数值以矩阵乘法C[4096,4096] A[4096,4096] × B[4096,4096]为例Naive每个线程算一个C[i,j]需要读 A 的一整行4096 个元素和 B 的一整列4096 个元素。每个元素做 4096 次乘加。算术强度 2×4096 / (2×4096×2字节) ≈ 0.5 FLOPs/ByteTiled把 C 切成 128×128 的块每块对应的 A 片段128×K和 B 片段K×128加载到 Shared Memory被块内所有线程复用。算术强度 2×128×128 / ((128128)×2字节) ≈ 64 FLOPs/Byte算术强度提升了 128 倍从 memory-bound 变成 compute-bound。案例解释前提设矩阵乘法CA×BCA×BCA×B其中AAA是M×KM×KM×KBBB是K×NK×NK×NCCC是M×NM×NM×N。这里MNK4096MNK4096MNK4096。计算每个C[i,j]C[i,j]C[i,j]需要K 次乘法KK 次加法所以每个元素2K2K2K次浮点运算FLOPs。整个矩阵乘法的总运算量总 FLOPsM×N×2K4096×4096×2×4096 \text{总 FLOPs} M \times N \times 2K 4096 \times 4096 \times 2 \times 4096总FLOPsM×N×2K4096×4096×2×4096但算术强度是单位字节内存访问所对应的运算量所以我们更关心的是每个元素计算时读取了多少数据。Naive 方法每个线程算一个C[i,j]C[i,j]C[i,j]计算量每个C[i,j]C[i,j]C[i,j]需要读取AAA的一整行KKK个元素和BBB的一整列KKK个元素。做KKK次乘法和KKK次加法共2K2K2K次 FLOPs。内存读取量每个C[i,j]C[i,j]C[i,j]读取AAA的行KKK个元素读取BBB的列KKK个元素总共2K2K2K个元素。假设每个元素是 2 字节fp16则总字节数 2K×24K2K×24K2K×24K字节。算术强度算术强度2K4K120.5 FLOPs/Byte \text{算术强度} \frac{2K}{4K} \frac{1}{2} 0.5\ \text{FLOPs/Byte}算术强度4K2K210.5FLOPs/Byte代入K4096K4096K40962×40962×4096×28192163840.5 \dfrac{2 \times 4096}{2 \times 4096 \times 2} \dfrac{8192}{16384} 0.52×4096×22×40961638481920.5Tiled 方法分块复用数据把CCC切成128×128128×128128×128的块。每个块的计算需要对应的AAA片段128×K128×K128×K对应的BBB片段K×128K×128K×128这些片段先加载到 Shared Memory然后被块内所有线程复用。每个 C 块的计算量块内有128×128128×128128×128个元素每个元素仍需2K2K2K次 FLOPsFLOPs128×128×2K FLOPs 128 \times 128 \times 2KFLOPs128×128×2K每个 C 块从全局内存读取的数据量A 片段128×K128×K128×K个元素B 片段K×128K×128K×128个元素总元素数 128K128K256K128K128K256K128K128K256K每个元素 2 字节总字节数 256K×2512K256K×2512K256K×2512K算术强度算术强度128×128×2K512K 算术强度 \frac{128 \times 128 \times 2K}{512K}算术强度512K128×128×2KK 在分子分母中约掉128×128×25123276851264 FLOPs/Byte \frac{128 \times 128 \times 2}{512} \frac{32768}{512} 64 \text{ FLOPs/Byte}512128×128×25123276864FLOPs/Byte原文中的公式2×128×128(128128)×23276851264 \frac{2 \times 128 \times 128}{(128 128) \times 2} \frac{32768}{512} 64(128128)×22×128×1285123276864这里分子 2×128×128 就是 128×128×2分母 (128128)×2 就是 256×2512。K 被省略是因为它在分子分母中消掉了。为什么提升了 128 倍Naive 算术强度0.5 FLOPs/ByteTiled 算术强度64 FLOPs/Byte提升倍数64/0.512864/0.5128根本原因在 Tiled 方法中每个 A 片段和 B 片段被加载到共享内存后被块内 128个线程对应 128 个不同的列或行复用了 128 次。一个 A 行元素比如A[i,k]A[i,k]A[i,k]原本在 Naive 中要被 N4096 个线程读取现在只被读取一次加载到共享内存然后被块内 128 个线程复用。同理B 元素也被复用了 128 次。因此全局内存流量减少了 128 倍而计算量不变所以算术强度提高了 128 倍。这直接把程序从 memory-bound受内存带宽限制变成了 compute-bound受计算能力限制大大提升了性能。H/W 维切分的通用公式概述这是 Tiling 在空间维度高度/宽度上切分时的核心公式。想象你要擦一面大玻璃窗输入图像你手里只有一块小抹布输出 tile。你不可能一次性擦完整面窗所以你把窗户分成几块一块一块地擦。但问题来了擦玻璃时抹布边缘会碰到隔壁区域所以相邻的两块之间必须有重叠否则交界处会擦不干净。HW 维切分就是干这件事把一个大图像切成若干小块分别计算同时精确算出每小块需要从原图里取多少数据、边界怎么补零、相邻块之间重叠多少。为什么需要切分卷积计算量很大比如一张 4096×4096 的图不可能一次性全部塞进 GPU 的显存或缓存里。所以必须切成小块一块一块算每块算完后把结果拼回完整输出切分后数据量小可以放进共享内存或高速缓存大幅加速约定符号含义H输入高度kh卷积核高度sh步长stridept顶部填充padding topd膨胀率dilationo0, o1输出 tile 的行范围 [o0, o1)tileOh输出 tile 的行数 o1 - o0公式输入切片起始行 in_start max(0, o0*sh - pt) 输入切片结束行(不含) in_end (o1 - 1)*sh - pt (kh - 1)*d 1 切片行数 in_span in_end - in_start 子算子顶部 pad pt max(0, pt - o0*sh) 子算子底部 pad pb max(0, in_end - H) 相邻 tile 的重叠行数 halo (kh - 1)*d 1 - sh逐条解释①in_start max(0, o0*sh - pt)这块输出从哪一行输入开始取输出第 o0 行的第一个窗口起点是o0*sh - pt。可能为负tile 顶部落在 pad 区所以 clamp 到 0。比如o00, sh2, pt1起点 0*2 - 1 -1但输入没有第 -1 行那是 pad所以从 0 开始即 max(0, -1) 0②in_end (o1-1)*sh - pt (kh-1)*d 1这块输出到哪一行输入结束输出最后一行o1-1的最后一个窗口终点。窗口是逐行滑动的行数只与 tileOh、kh、d 有关。它对应的窗口起点是(o1-1)*sh - pt窗口本身有(kh-1)*d 1行卷积核高度所以结束位置 起点 窗口高度③pt max(0, pt - o0*sh)切出来的这一小块顶部还需要补零吗原来整张图顶部补了 pt 行零但切出来的这块可能已经跳过了那些零因为 o0*sh 已经往下走了如果切片的起点在原图第 0 行或之后就不需要再补零了通俗理解本来窗户顶部贴了胶带pad我现在从中间某处开始擦胶带已经被跳过了不需要再贴。④pb max(0, in_end - H)切出来的这一小块底部还需要补零吗slice 底部超出图像的部分要补回 pb’否则子卷积的结果高度会变少。in_end是需要的输入结束行如果超过了原图高度H说明超出了图像底部超出的部分要补零补多少就是in_end - H通俗理解我要取到第 10 行但图只有 9 行那第 10 行就得用零补上。⑤halo (kh-1)*d 1 - sh相邻两块之间重叠多少行一个窗口需要(kh-1)*d 1行输入每多一个输出行输入只多走sh行多出来的(kh-1)*d 1 - sh行就是被下一块复用的部分通俗理解我擦完这块玻璃下一块要和我重叠一点不然交界处擦不干净。重叠多少呢就是halo。详细数值例子参数H9, kh3, sh2, pt1, d1输出OH4tileOh2。输入行: -1 0 1 2 3 4 5 6 7 8 9 [pad] [x0] [x1] [x2] [x3] [x4] [x5] [x6] [x7] [x8] [pad] 输出 o0 的窗口: [pad] [x0] [x1] 输出 o1 的窗口: [x1] [x2] [x3] 输出 o2 的窗口: [x3] [x4] [x5] 输出 o3 的窗口: [x5] [x6] [x7]tile Ao00, o12in_start max(0, 0*2 - 1) max(0, -1) 0 in_end (2-1)*2 - 1 (3-1)*1 1 2 - 1 2 1 4 in_span 4 - 0 4 pt max(0, 1 - 0) 1 pb max(0, 4 - 9) 0 切输入 [x0..x3]子算子顶部 pad 1tile Bo02, o14in_start max(0, 2*2 - 1) max(0, 3) 3 in_end (4-1)*2 - 1 (3-1)*1 1 6 - 1 2 1 8 in_span 8 - 3 5 pt max(0, 1 - 4) 0 pb max(0, 8 - 9) 0 切输入 [x3..x7]子算子顶部 pad 0halohalo (3-1)*1 1 - 2 2 1 - 2 1 行即x3被 tile A 和 tile B 同时使用这就是 halo。为什么教学工程先切 batch 维batch 维切分各片互不相交天然零 halo// 切分前 %0 top.Conv2D %in : tensor2x3x16x16xf32 - tensor2x8x16x16xf32 // 切分后tile-size1 %e tensor.empty : tensor2x8x16x16xf32 %s0 tensor.extract_slice %in[0,0,0,0][1,3,16,16][1,1,1,1] %c0 top.Conv2D %s0 : - tensor1x8x16x16xf32 %a0 tensor.insert_slice %c0 into %e[0,0,0,0] %s1 tensor.extract_slice %in[1,0,0,0][1,3,16,16][1,1,1,1] %c1 top.Conv2D %s1 : - tensor1x8x16x16xf32 %d tensor.insert_slice %c1 into %a0[1,0,0,0]batch 0 和 batch 1 完全不重叠不需要处理 halo也不需要修正 pads。Batch 0: 图像 A → 输出 A Batch 1: 图像 B → 输出 B没有重叠halo 0不需要修正 pad不需要处理边界多级 Tiling 策略多级 Tiling 是把一个大矩阵运算逐层分解使每一级的数据都能驻留在对应的存储层级中。以 GEMMC A × B为例CUTLASS 采用三级层次Thread Block Level最外层 → 输出 C 分成 BLOCK_M × BLOCK_N 的块每块由一个 thread block 计算 → 数据从 HBM 加载到 Shared Memory Warp Level中间层 → 一个 thread block 内进一步把 BLOCK_M × BLOCK_N 分给多个 warp → 每个 warp 负责 WARP_M × WARP_N 的子块 → 数据从 Shared Memory 读入 Register File Register Level / MMA最内层 → 最内层映射到 Tensor Core 的 MMA 指令 → 每条 MMA 处理 16×8×16 或 16×16×16 的小矩阵块 → 数据以 fragment 形式分散在 warp 中 32 个线程的 register 中具体数值以 A100 上C[4096,4096] A[4096,4096] × B[4096,4096]FP16为例层级Tile 大小存储位置带宽延迟Thread Block128×128Shared Memory~19 TB/s~30 cyclesWarp64×64Shared Memory → Registeron-chip~1 cycleMMA16×16×16Registeron-chip~1 cycleThread Block 级C 分成128×128的块共32×32 1024个 thread block。每个 thread block 沿 K 维迭代每次加载 128×32 的 A 片段和32×128的 B 片段到 Shared Memory。Warp 级128×128的 thread block 分给 4 个 warp2×2 布局每个 warp 处理64×64的子块。Register/MMA 级每个 warp 用m16n8k16的 MMA 指令处理 16×8×16 的矩阵乘法。数据以 fragment 形式分散在 32 个线程的 register 中。嵌套关系Thread Block Tile: 128×128 └── Warp Tile: 64×64 └── MMA Tile: 16×16×16每一级的数据都驻留在对应的存储层级带宽逐级提升约 10 倍。内存层次优化技术选定 tile size 后还需要一系列技术确保数据在各级存储之间高效流动。主要有五种Shared Memory Staging共享内存暂存最基本的模式外层循环沿 K 维步进每步把一个 tile 从 HBM 加载到 Shared Memory内层循环在 Shared Memory 中执行计算。__shared__ half A_smem[BLOCK_M][BLOCK_K];__shared__ half B_smem[BLOCK_K][BLOCK_N];for(intk0;kK;kBLOCK_K){load_tile_A(A,A_smem,k);// HBM → Shared Memoryload_tile_B(B,B_smem,k);__syncthreads();// 等所有线程完成 loadcompute_tile(A_smem,B_smem,C_reg);// Shared Memory → Register__syncthreads();// 等所有线程完成 compute}cp.async异步复制Ampere 之前HBM → Shared Memory 的路径是HBM → global load → Register → shared store → Shared Memory要占用 register 作为中转。Ampere 引入 cp.async允许数据直接从 HBM 复制到 Shared Memory不经过 registerHBM → cp.async → Shared Memory好处有三减少 register 压力允许 load 和 compute 重叠提高 SM 资源利用率Double Buffering / Multi-Stage Pipelining如果 load 和 compute 完全串行一半时间 Tensor Core 在等数据。Double Buffering 解决这个问题Iteration i: Buffer A: Compute(Tile[i]) ← Tensor Core busy Buffer B: Load(Tile[i1]) ← Memory pipeline busy (cp.async) Iteration i1: Buffer B: Compute(Tile[i1]) ← swap Buffer A: Load(Tile[i2]) ← swap可以推广到 N-stage pipelineN 3, 4, 5每多一级需要多分配一份 Shared Memory buffer。SMEM usage tile_size × elem_size × num_stages更多 stage 意味着更好的 latency hiding但也意味着更多 Shared Memory 消耗可能导致 occupancy 下降。Memory Coalescing合并访问一个 warp 中的 32 个线程应该访问连续的内存地址这样硬件可以将 32 次请求合并为 1 次 128-byte transaction。// Coalesced: thread i reads element idata[threadIdx.x]// ✓ One 128B transaction// Non-coalesced: thread i reads element i*stridedata[threadIdx.x*stride]// ✗ Multiple transactions非合并访问最坏情况下32 个线程产生 32 次独立的内存事务性能下降 10~32 倍。Bank Conflict 与 SwizzlingShared Memory 由 32 个 bank 组成每个 bank 宽 4 字节。如果多个线程访问同一个 bank必须串行化。访问模式Bank 分布冲突度性能影响stride10,1,2,…,311-way无冲突1×stride20,2,4,…,30,0,2,…2-way2× 慢stride320,0,0,…,032-way全冲突32× 慢Swizzling是消除 bank conflict 的标准技术对地址做 XOR 运算重新映射 bank 分配。swizzled_bank (row ⊕ col) mod 32Tile Size 选择的约束分析Tile size 的选择受多个硬件约束限制形成可行性区域。约束一Shared Memory 容量SMEM (BLOCK_M × BLOCK_K BLOCK_K × BLOCK_N) × elem_size × num_stages以BLOCK_M BLOCK_N 128, BLOCK_K 32, FP16, 2-stage为例SMEM (128×32 32×128) × 2 × 2 32,768 bytes 32 KBA100 每个 SM 有约 164 KB Shared Memory。32 KB 的 tile 理论上可以容纳 5 个 thread block164/32 5。如果增大到BLOCK_M BLOCK_N 256, BLOCK_K 64, 3-stageSMEM (256×64 64×256) × 2 × 3 196,608 bytes 192 KB超过 A100 的 164 KB 限制 → 不可行。约束二Register 压力每个线程需要持有 accumulatorregs_per_thread ≈ (BLOCK_M × BLOCK_N) / threads_per_block overhead以BLOCK_M BLOCK_N 128, threads_per_block 1024为例regs_per_thread ≈ (128×128) / 1024 32 16 32 48NVIDIA GPU 每个线程最多使用 255 个 32-bit register。如果超过编译器会把 register 溢出到 local memory实际上是 HBM导致性能灾难。约束三OccupancyOccupancy SM 上实际活跃的 warp 数 / SM 支持的最大 warp 数。更大的 tile → 更多 Shared Memory 和 register → 每个 SM 能运行的 thread block 更少 → occupancy 下降。低 occupancy 意味着当一些 warp 在等内存时没有足够的其他 warp 填补空闲的计算周期。一般来说occupancy 低于 25% 会导致明显的性能下降。但 occupancy 也不是越高越好——更大的 tile 意味着更高的数据复用率。这是一个经典的 occupancy vs. data reuse 权衡。自动调优实践中往往通过 autotuning 寻找最优 tile size① 枚举所有可行的 tile 配置组合 ② 过滤掉违反硬件约束的配置 ③ 在目标硬件上对每个候选配置进行基准测试 ④ 选择性能最高的配置Triton 的autotunedecorator 和 CUTLASS 的 profiling 工具都是这个思路。经验法则情况调整起始点BLOCK_M BLOCK_N 128, BLOCK_K 32Shared Memory 有余量增大 BLOCK_K 或增加 pipeline stageOccupancy 过低减小 tile sizeRegister 溢出减小 BLOCK_M × BLOCK_N总结问题答案Tiling 是什么把大计算切成小块让数据装进片上高速存储H/W 维切分公式in_start max(0, o0*sh - pt);in_end (o1-1)*sh - pt (kh-1)*d 1;pt max(0, pt - o0*sh);pb max(0, in_end - H);halo (kh-1)*d 1 - sh多级 TilingThread Block → Warp → Register/MMA三级映射到 HBM → Shared Memory → Register内存层次优化Shared Memory Staging、cp.async、Double Buffering、Memory Coalescing、SwizzlingTile Size 约束Shared Memory 容量、Register 压力、Occupancy 三角约束一句话Tiling 是连接高层算法优化和底层硬件执行的桥梁。它把大矩阵切成小块让数据在 HBM → Shared Memory → Register 的层次结构中逐级驻留通过 Double Buffering 隐藏延迟、Swizzling 消除 bank conflict、约束分析选择最优 tile size最终把 memory-bound 问题转化为 compute-bound。
Tiling（分块）
概述Tiling分块把一个大计算或大张量切成许多小块每块单独计算让数据能装进片上高速存储。例子餐厅比喻Naive不分块 从仓库HBM拿一袋米 → 跑到厨房SM→ 用一勺 → 跑回仓库放好 → 再拿一袋 → 再跑 → 再用一勺 → ... 每次用一勺米都要跑一趟仓库累死。 Tiled分块 从仓库一次搬一整袋米Tile到厨房台面Shared Memory → 在台面上反复用Compute → 用完了再搬下一袋 跑仓库的次数大幅减少。具体数值以矩阵乘法C[4096,4096] A[4096,4096] × B[4096,4096]为例Naive每个线程算一个C[i,j]需要读 A 的一整行4096 个元素和 B 的一整列4096 个元素。每个元素做 4096 次乘加。算术强度 2×4096 / (2×4096×2字节) ≈ 0.5 FLOPs/ByteTiled把 C 切成 128×128 的块每块对应的 A 片段128×K和 B 片段K×128加载到 Shared Memory被块内所有线程复用。算术强度 2×128×128 / ((128128)×2字节) ≈ 64 FLOPs/Byte算术强度提升了 128 倍从 memory-bound 变成 compute-bound。案例解释前提设矩阵乘法CA×BCA×BCA×B其中AAA是M×KM×KM×KBBB是K×NK×NK×NCCC是M×NM×NM×N。这里MNK4096MNK4096MNK4096。计算每个C[i,j]C[i,j]C[i,j]需要K 次乘法KK 次加法所以每个元素2K2K2K次浮点运算FLOPs。整个矩阵乘法的总运算量总 FLOPsM×N×2K4096×4096×2×4096 \text{总 FLOPs} M \times N \times 2K 4096 \times 4096 \times 2 \times 4096总FLOPsM×N×2K4096×4096×2×4096但算术强度是单位字节内存访问所对应的运算量所以我们更关心的是每个元素计算时读取了多少数据。Naive 方法每个线程算一个C[i,j]C[i,j]C[i,j]计算量每个C[i,j]C[i,j]C[i,j]需要读取AAA的一整行KKK个元素和BBB的一整列KKK个元素。做KKK次乘法和KKK次加法共2K2K2K次 FLOPs。内存读取量每个C[i,j]C[i,j]C[i,j]读取AAA的行KKK个元素读取BBB的列KKK个元素总共2K2K2K个元素。假设每个元素是 2 字节fp16则总字节数 2K×24K2K×24K2K×24K字节。算术强度算术强度2K4K120.5 FLOPs/Byte \text{算术强度} \frac{2K}{4K} \frac{1}{2} 0.5\ \text{FLOPs/Byte}算术强度4K2K210.5FLOPs/Byte代入K4096K4096K40962×40962×4096×28192163840.5 \dfrac{2 \times 4096}{2 \times 4096 \times 2} \dfrac{8192}{16384} 0.52×4096×22×40961638481920.5Tiled 方法分块复用数据把CCC切成128×128128×128128×128的块。每个块的计算需要对应的AAA片段128×K128×K128×K对应的BBB片段K×128K×128K×128这些片段先加载到 Shared Memory然后被块内所有线程复用。每个 C 块的计算量块内有128×128128×128128×128个元素每个元素仍需2K2K2K次 FLOPsFLOPs128×128×2K FLOPs 128 \times 128 \times 2KFLOPs128×128×2K每个 C 块从全局内存读取的数据量A 片段128×K128×K128×K个元素B 片段K×128K×128K×128个元素总元素数 128K128K256K128K128K256K128K128K256K每个元素 2 字节总字节数 256K×2512K256K×2512K256K×2512K算术强度算术强度128×128×2K512K 算术强度 \frac{128 \times 128 \times 2K}{512K}算术强度512K128×128×2KK 在分子分母中约掉128×128×25123276851264 FLOPs/Byte \frac{128 \times 128 \times 2}{512} \frac{32768}{512} 64 \text{ FLOPs/Byte}512128×128×25123276864FLOPs/Byte原文中的公式2×128×128(128128)×23276851264 \frac{2 \times 128 \times 128}{(128 128) \times 2} \frac{32768}{512} 64(128128)×22×128×1285123276864这里分子 2×128×128 就是 128×128×2分母 (128128)×2 就是 256×2512。K 被省略是因为它在分子分母中消掉了。为什么提升了 128 倍Naive 算术强度0.5 FLOPs/ByteTiled 算术强度64 FLOPs/Byte提升倍数64/0.512864/0.5128根本原因在 Tiled 方法中每个 A 片段和 B 片段被加载到共享内存后被块内 128个线程对应 128 个不同的列或行复用了 128 次。一个 A 行元素比如A[i,k]A[i,k]A[i,k]原本在 Naive 中要被 N4096 个线程读取现在只被读取一次加载到共享内存然后被块内 128 个线程复用。同理B 元素也被复用了 128 次。因此全局内存流量减少了 128 倍而计算量不变所以算术强度提高了 128 倍。这直接把程序从 memory-bound受内存带宽限制变成了 compute-bound受计算能力限制大大提升了性能。H/W 维切分的通用公式概述这是 Tiling 在空间维度高度/宽度上切分时的核心公式。想象你要擦一面大玻璃窗输入图像你手里只有一块小抹布输出 tile。你不可能一次性擦完整面窗所以你把窗户分成几块一块一块地擦。但问题来了擦玻璃时抹布边缘会碰到隔壁区域所以相邻的两块之间必须有重叠否则交界处会擦不干净。HW 维切分就是干这件事把一个大图像切成若干小块分别计算同时精确算出每小块需要从原图里取多少数据、边界怎么补零、相邻块之间重叠多少。为什么需要切分卷积计算量很大比如一张 4096×4096 的图不可能一次性全部塞进 GPU 的显存或缓存里。所以必须切成小块一块一块算每块算完后把结果拼回完整输出切分后数据量小可以放进共享内存或高速缓存大幅加速约定符号含义H输入高度kh卷积核高度sh步长stridept顶部填充padding topd膨胀率dilationo0, o1输出 tile 的行范围 [o0, o1)tileOh输出 tile 的行数 o1 - o0公式输入切片起始行 in_start max(0, o0*sh - pt) 输入切片结束行(不含) in_end (o1 - 1)*sh - pt (kh - 1)*d 1 切片行数 in_span in_end - in_start 子算子顶部 pad pt max(0, pt - o0*sh) 子算子底部 pad pb max(0, in_end - H) 相邻 tile 的重叠行数 halo (kh - 1)*d 1 - sh逐条解释①in_start max(0, o0*sh - pt)这块输出从哪一行输入开始取输出第 o0 行的第一个窗口起点是o0*sh - pt。可能为负tile 顶部落在 pad 区所以 clamp 到 0。比如o00, sh2, pt1起点 0*2 - 1 -1但输入没有第 -1 行那是 pad所以从 0 开始即 max(0, -1) 0②in_end (o1-1)*sh - pt (kh-1)*d 1这块输出到哪一行输入结束输出最后一行o1-1的最后一个窗口终点。窗口是逐行滑动的行数只与 tileOh、kh、d 有关。它对应的窗口起点是(o1-1)*sh - pt窗口本身有(kh-1)*d 1行卷积核高度所以结束位置 起点 窗口高度③pt max(0, pt - o0*sh)切出来的这一小块顶部还需要补零吗原来整张图顶部补了 pt 行零但切出来的这块可能已经跳过了那些零因为 o0*sh 已经往下走了如果切片的起点在原图第 0 行或之后就不需要再补零了通俗理解本来窗户顶部贴了胶带pad我现在从中间某处开始擦胶带已经被跳过了不需要再贴。④pb max(0, in_end - H)切出来的这一小块底部还需要补零吗slice 底部超出图像的部分要补回 pb’否则子卷积的结果高度会变少。in_end是需要的输入结束行如果超过了原图高度H说明超出了图像底部超出的部分要补零补多少就是in_end - H通俗理解我要取到第 10 行但图只有 9 行那第 10 行就得用零补上。⑤halo (kh-1)*d 1 - sh相邻两块之间重叠多少行一个窗口需要(kh-1)*d 1行输入每多一个输出行输入只多走sh行多出来的(kh-1)*d 1 - sh行就是被下一块复用的部分通俗理解我擦完这块玻璃下一块要和我重叠一点不然交界处擦不干净。重叠多少呢就是halo。详细数值例子参数H9, kh3, sh2, pt1, d1输出OH4tileOh2。输入行: -1 0 1 2 3 4 5 6 7 8 9 [pad] [x0] [x1] [x2] [x3] [x4] [x5] [x6] [x7] [x8] [pad] 输出 o0 的窗口: [pad] [x0] [x1] 输出 o1 的窗口: [x1] [x2] [x3] 输出 o2 的窗口: [x3] [x4] [x5] 输出 o3 的窗口: [x5] [x6] [x7]tile Ao00, o12in_start max(0, 0*2 - 1) max(0, -1) 0 in_end (2-1)*2 - 1 (3-1)*1 1 2 - 1 2 1 4 in_span 4 - 0 4 pt max(0, 1 - 0) 1 pb max(0, 4 - 9) 0 切输入 [x0..x3]子算子顶部 pad 1tile Bo02, o14in_start max(0, 2*2 - 1) max(0, 3) 3 in_end (4-1)*2 - 1 (3-1)*1 1 6 - 1 2 1 8 in_span 8 - 3 5 pt max(0, 1 - 4) 0 pb max(0, 8 - 9) 0 切输入 [x3..x7]子算子顶部 pad 0halohalo (3-1)*1 1 - 2 2 1 - 2 1 行即x3被 tile A 和 tile B 同时使用这就是 halo。为什么教学工程先切 batch 维batch 维切分各片互不相交天然零 halo// 切分前 %0 top.Conv2D %in : tensor2x3x16x16xf32 - tensor2x8x16x16xf32 // 切分后tile-size1 %e tensor.empty : tensor2x8x16x16xf32 %s0 tensor.extract_slice %in[0,0,0,0][1,3,16,16][1,1,1,1] %c0 top.Conv2D %s0 : - tensor1x8x16x16xf32 %a0 tensor.insert_slice %c0 into %e[0,0,0,0] %s1 tensor.extract_slice %in[1,0,0,0][1,3,16,16][1,1,1,1] %c1 top.Conv2D %s1 : - tensor1x8x16x16xf32 %d tensor.insert_slice %c1 into %a0[1,0,0,0]batch 0 和 batch 1 完全不重叠不需要处理 halo也不需要修正 pads。Batch 0: 图像 A → 输出 A Batch 1: 图像 B → 输出 B没有重叠halo 0不需要修正 pad不需要处理边界多级 Tiling 策略多级 Tiling 是把一个大矩阵运算逐层分解使每一级的数据都能驻留在对应的存储层级中。以 GEMMC A × B为例CUTLASS 采用三级层次Thread Block Level最外层 → 输出 C 分成 BLOCK_M × BLOCK_N 的块每块由一个 thread block 计算 → 数据从 HBM 加载到 Shared Memory Warp Level中间层 → 一个 thread block 内进一步把 BLOCK_M × BLOCK_N 分给多个 warp → 每个 warp 负责 WARP_M × WARP_N 的子块 → 数据从 Shared Memory 读入 Register File Register Level / MMA最内层 → 最内层映射到 Tensor Core 的 MMA 指令 → 每条 MMA 处理 16×8×16 或 16×16×16 的小矩阵块 → 数据以 fragment 形式分散在 warp 中 32 个线程的 register 中具体数值以 A100 上C[4096,4096] A[4096,4096] × B[4096,4096]FP16为例层级Tile 大小存储位置带宽延迟Thread Block128×128Shared Memory~19 TB/s~30 cyclesWarp64×64Shared Memory → Registeron-chip~1 cycleMMA16×16×16Registeron-chip~1 cycleThread Block 级C 分成128×128的块共32×32 1024个 thread block。每个 thread block 沿 K 维迭代每次加载 128×32 的 A 片段和32×128的 B 片段到 Shared Memory。Warp 级128×128的 thread block 分给 4 个 warp2×2 布局每个 warp 处理64×64的子块。Register/MMA 级每个 warp 用m16n8k16的 MMA 指令处理 16×8×16 的矩阵乘法。数据以 fragment 形式分散在 32 个线程的 register 中。嵌套关系Thread Block Tile: 128×128 └── Warp Tile: 64×64 └── MMA Tile: 16×16×16每一级的数据都驻留在对应的存储层级带宽逐级提升约 10 倍。内存层次优化技术选定 tile size 后还需要一系列技术确保数据在各级存储之间高效流动。主要有五种Shared Memory Staging共享内存暂存最基本的模式外层循环沿 K 维步进每步把一个 tile 从 HBM 加载到 Shared Memory内层循环在 Shared Memory 中执行计算。__shared__ half A_smem[BLOCK_M][BLOCK_K];__shared__ half B_smem[BLOCK_K][BLOCK_N];for(intk0;kK;kBLOCK_K){load_tile_A(A,A_smem,k);// HBM → Shared Memoryload_tile_B(B,B_smem,k);__syncthreads();// 等所有线程完成 loadcompute_tile(A_smem,B_smem,C_reg);// Shared Memory → Register__syncthreads();// 等所有线程完成 compute}cp.async异步复制Ampere 之前HBM → Shared Memory 的路径是HBM → global load → Register → shared store → Shared Memory要占用 register 作为中转。Ampere 引入 cp.async允许数据直接从 HBM 复制到 Shared Memory不经过 registerHBM → cp.async → Shared Memory好处有三减少 register 压力允许 load 和 compute 重叠提高 SM 资源利用率Double Buffering / Multi-Stage Pipelining如果 load 和 compute 完全串行一半时间 Tensor Core 在等数据。Double Buffering 解决这个问题Iteration i: Buffer A: Compute(Tile[i]) ← Tensor Core busy Buffer B: Load(Tile[i1]) ← Memory pipeline busy (cp.async) Iteration i1: Buffer B: Compute(Tile[i1]) ← swap Buffer A: Load(Tile[i2]) ← swap可以推广到 N-stage pipelineN 3, 4, 5每多一级需要多分配一份 Shared Memory buffer。SMEM usage tile_size × elem_size × num_stages更多 stage 意味着更好的 latency hiding但也意味着更多 Shared Memory 消耗可能导致 occupancy 下降。Memory Coalescing合并访问一个 warp 中的 32 个线程应该访问连续的内存地址这样硬件可以将 32 次请求合并为 1 次 128-byte transaction。// Coalesced: thread i reads element idata[threadIdx.x]// ✓ One 128B transaction// Non-coalesced: thread i reads element i*stridedata[threadIdx.x*stride]// ✗ Multiple transactions非合并访问最坏情况下32 个线程产生 32 次独立的内存事务性能下降 10~32 倍。Bank Conflict 与 SwizzlingShared Memory 由 32 个 bank 组成每个 bank 宽 4 字节。如果多个线程访问同一个 bank必须串行化。访问模式Bank 分布冲突度性能影响stride10,1,2,…,311-way无冲突1×stride20,2,4,…,30,0,2,…2-way2× 慢stride320,0,0,…,032-way全冲突32× 慢Swizzling是消除 bank conflict 的标准技术对地址做 XOR 运算重新映射 bank 分配。swizzled_bank (row ⊕ col) mod 32Tile Size 选择的约束分析Tile size 的选择受多个硬件约束限制形成可行性区域。约束一Shared Memory 容量SMEM (BLOCK_M × BLOCK_K BLOCK_K × BLOCK_N) × elem_size × num_stages以BLOCK_M BLOCK_N 128, BLOCK_K 32, FP16, 2-stage为例SMEM (128×32 32×128) × 2 × 2 32,768 bytes 32 KBA100 每个 SM 有约 164 KB Shared Memory。32 KB 的 tile 理论上可以容纳 5 个 thread block164/32 5。如果增大到BLOCK_M BLOCK_N 256, BLOCK_K 64, 3-stageSMEM (256×64 64×256) × 2 × 3 196,608 bytes 192 KB超过 A100 的 164 KB 限制 → 不可行。约束二Register 压力每个线程需要持有 accumulatorregs_per_thread ≈ (BLOCK_M × BLOCK_N) / threads_per_block overhead以BLOCK_M BLOCK_N 128, threads_per_block 1024为例regs_per_thread ≈ (128×128) / 1024 32 16 32 48NVIDIA GPU 每个线程最多使用 255 个 32-bit register。如果超过编译器会把 register 溢出到 local memory实际上是 HBM导致性能灾难。约束三OccupancyOccupancy SM 上实际活跃的 warp 数 / SM 支持的最大 warp 数。更大的 tile → 更多 Shared Memory 和 register → 每个 SM 能运行的 thread block 更少 → occupancy 下降。低 occupancy 意味着当一些 warp 在等内存时没有足够的其他 warp 填补空闲的计算周期。一般来说occupancy 低于 25% 会导致明显的性能下降。但 occupancy 也不是越高越好——更大的 tile 意味着更高的数据复用率。这是一个经典的 occupancy vs. data reuse 权衡。自动调优实践中往往通过 autotuning 寻找最优 tile size① 枚举所有可行的 tile 配置组合 ② 过滤掉违反硬件约束的配置 ③ 在目标硬件上对每个候选配置进行基准测试 ④ 选择性能最高的配置Triton 的autotunedecorator 和 CUTLASS 的 profiling 工具都是这个思路。经验法则情况调整起始点BLOCK_M BLOCK_N 128, BLOCK_K 32Shared Memory 有余量增大 BLOCK_K 或增加 pipeline stageOccupancy 过低减小 tile sizeRegister 溢出减小 BLOCK_M × BLOCK_N总结问题答案Tiling 是什么把大计算切成小块让数据装进片上高速存储H/W 维切分公式in_start max(0, o0*sh - pt);in_end (o1-1)*sh - pt (kh-1)*d 1;pt max(0, pt - o0*sh);pb max(0, in_end - H);halo (kh-1)*d 1 - sh多级 TilingThread Block → Warp → Register/MMA三级映射到 HBM → Shared Memory → Register内存层次优化Shared Memory Staging、cp.async、Double Buffering、Memory Coalescing、SwizzlingTile Size 约束Shared Memory 容量、Register 压力、Occupancy 三角约束一句话Tiling 是连接高层算法优化和底层硬件执行的桥梁。它把大矩阵切成小块让数据在 HBM → Shared Memory → Register 的层次结构中逐级驻留通过 Double Buffering 隐藏延迟、Swizzling 消除 bank conflict、约束分析选择最优 tile size最终把 memory-bound 问题转化为 compute-bound。
