为什么在 SIMT 之外还需要 Tile

SIMT 让程序员明确安排 thread、warp、shared memory 和同步，控制力很强，但矩阵/张量 kernel 的实现容易绑死在某代硬件细节上。CUDA Tile 把 block 内的并行表达为对多维 tile 值的整体运算，由编译器决定线程如何协作，并选择 TMA、Tensor Core 等实现。

SIMT 视角                              Tile 视角
thread 0..N 各自 load/compute/store    load 一个 [M,N] tile
程序员组织 warp 和 shared memory       对 tile 做 +、reduce、matmul
程序员安排 barrier/pipeline            编译器降低到线程、搬运和矩阵指令

Tile 不是 SIMT 的语法糖，也不是要取代所有线程级 kernel。它是另一种设备编程模型，特别适合规则的块状张量计算。

本章按 CUDA Toolkit 13.4 文档整理。CUDA Tile Python 通过 cuda.tile 提供；CUDA Tile C++ 从 Toolkit 13.3 起提供，接口仍在快速发展。生产采用前必须锁定 Toolkit、Driver、目标架构并重新核对当前 API。

两种模型如何共存

同一翻译单元可以包含 Tile 与 SIMT kernel，但当前两类 device function 不能任意互相调用。host 端可以按工作负载分别 launch，两者读写相同的标准 CUDA device allocations。

核心心智模型

以二维矩阵为例：

Global tensor A[M,N]

                 tile program grid
          x=0          x=1          x=2
       ┌────────┬────────────┬────────────┐
y=0    │ tile   │ tile       │ tile       │
       │ [BM,BN]│ [BM,BN]    │ [BM,BN]    │
       ├────────┼────────────┼────────────┤
y=1    │ tile   │ tile       │ masked edge│
       └────────┴────────────┴────────────┘

每个 program instance：
bid → 定位分区 → load tile → tile 运算 → store tile

Tile 的维度和 shape 通常是编译期信息，并受 API 的合法 shape 规则约束；边界不整除时使用 masked load/store，而不是让每个显式线程写 if (i<n)。

CUDA Tile C++ 最小例子

下面采用 13.4 官方参考中的主要结构，展示 1024 元素分区的向量加法：

#include "cuda_tile.h"

namespace ct = cuda::tiles;
using namespace ct::literals;

__tile_global__
void vector_add(const float* __restrict__ a,
                const float* __restrict__ b,
                float* __restrict__ out,
                std::size_t n) {
    a = assume_aligned(a, 16_ic);
    b = assume_aligned(b, 16_ic);
    out = assume_aligned(out, 16_ic);

    auto partition = ct::bid().x;
    auto va = ct::partition_view{
        ct::tensor_span{a, ct::extents{n}}, ct::shape{1024_ic}};
    auto vb = ct::partition_view{
        ct::tensor_span{b, ct::extents{n}}, ct::shape{1024_ic}};
    auto vo = ct::partition_view{
        ct::tensor_span{out, ct::extents{n}}, ct::shape{1024_ic}};

    auto ta = va.load_masked(partition);
    auto tb = vb.load_masked(partition);
    vo.store_masked(ta + tb, partition);
}

当前 Toolkit 的编译和启动形态：

nvcc -std=c++20 -enable-tile -arch=sm_80 app.cu -o app

std::size_t programs = (n + 1023) / 1024;
vector_add<<<programs, 1>>>(a, b, out, n);

这里的 1 不表示只用一个硬件线程完成 1024 次工作。Tile 编译器根据 tile 运算决定 block 内的实际并行实现。不要把 SIMT 的 threadIdx.x 推理直接套到 Tile kernel。

Tensor、span、view 与 tile value

四层对象各司其职：

device pointer + extents
        ↓
tensor_span：描述全局张量及形状
        ↓
partition/slice view：描述如何把全局张量分区
        ↓ load
tile value：片上值，可参与逐元素、归约、重排、矩阵运算
        ↓ store
global tensor

• span/view 通常只是描述符，不等于已经搬运数据；
• load/store 才建立全局内存流量；
• tile value 的 shape 是类型和优化的重要组成；
• masked 版本负责不完整边界，但内部 tile 仍保持规则 shape；
• 对齐承诺必须与真实 pointer/stride 一致。

Tile 算法模式

逐元素与广播

x tile [M,N] ─┐
              ├─ y = relu(x * scale + bias) → store
bias [N] ─────┘           ↑ broadcast

这类操作适合融合，减少中间数组和 launch。但融合后 live tile 太多也会增加寄存器/片上存储压力，需要测量。

归约

Tile reduction 显式指定轴：

input [M,N]
├─ reduce axis=1 → [M]  行归约
└─ reduce axis=0 → [N]  列归约

softmax 的稳定形式可以表达为：

$$m_i=\max_j x_{ij},\quad s_i=\sum_j e^{x_{ij}-m_i},\quad y_{ij}=\frac{e^{x_{ij}-m_i}}{s_i}$$

编写时仍要关注归约 dtype、NaN 传播、空维度和极端输入。

矩阵乘法

A tile [BM,BK] ─┐
                ├─ matmul / mma → accumulator [BM,BN]
B tile [BK,BN] ─┘
                         ↓
             epilogue: alpha, bias, activation, cast

Tile 模型的主要价值之一，是让编译器将规则矩阵操作映射到目标架构支持的 Tensor Core 与数据移动路径。性能仍取决于 shape、dtype、布局和边界；高层表达不会让任意尺寸自动达到库峰值。

Tile 与 TMA

支持的硬件和布局上，Tile load/store 可降低到 Tensor Memory Accelerator。TMA 擅长大块、多维 global↔shared 搬运，地址生成工作由专用硬件承担，并可与计算重叠。

global tensor
    │ descriptor: base + shape + stride + tile box
    ▼
TMA bulk transfer ── arrive/barrier ── tile compute
    ▲                                      │
    └──────────── next tile prefetch ──────┘

Tile 程序员描述“要哪个 tile”；编译器决定是否和如何使用 TMA。判断是否真的生成期望路径，仍需看编译报告、SASS/中间表示和 profiler，而不是从源码抽象层推断。

Tile Python 的位置

CUDA Tile Python 通过 cuda.tile 提供，概念与 C++ 一致：block id、global array、tile load/store、逐元素和矩阵/归约 primitive。其价值是用 Python 表达 kernel，同时保留编译期 tile shape。

import cuda.tile as ct
import cupy

TILE_SIZE = 16

@ct.kernel
def vector_add_kernel(a, b, result):
    block_id = ct.bid(0)
    a_tile = ct.load(a, index=(block_id,), shape=(TILE_SIZE,))
    b_tile = ct.load(b, index=(block_id,), shape=(TILE_SIZE,))
    ct.store(result, index=(block_id,), tile=a_tile + b_tile)

def vector_add(a: cupy.ndarray, b: cupy.ndarray, result: cupy.ndarray):
    grid = (ct.cdiv(a.shape[0], TILE_SIZE), 1, 1)
    ct.launch(cupy.cuda.get_current_stream(), grid,
              vector_add_kernel, (a, b, result))

这是当前官方 quickstart 的核心结构。实际代码还应验证输入 shape/dtype/设备，并按当前 cuTile 版本处理非整除边界。不要把其他 tile DSL 的 API 名直接当作 cuda.tile API；不同 DSL 的思想相似，但 decorator、shape 系统、布局描述和编译工具链并不等价。

如何从 SIMT 迁移思路

不要逐行翻译；先提升算法层级：

适合优先迁移的 kernel：GEMM-like、卷积 lowering、attention 子块、规则 reduction、逐元素融合。图遍历、复杂队列、细粒度原子协议和高度数据依赖控制流通常仍适合 SIMT。

性能调优维度

1. Program tile shape：影响并行实例数、边界浪费与数据复用。
2. 数据类型：决定 Tensor Core 路径、累加精度和带宽。
3. 布局与 stride：决定连续性、TMA 可用性和转置成本。
4. 融合边界：减少流量，但增加 live values 与编译复杂度。
5. 目标架构：同一表达可能在不同 GPU 上选择不同实现。
6. 问题尺寸分布：只调一个方阵尺寸会产生错误结论。

基准必须与 cuBLAS/cuBLASLt、CUTLASS 或已有 SIMT 实现比较。Tile 的目标是提高生产力和可移植性能，不保证每个手写表达都优于高度调优的专用库。

当前边界与部署注意

• CUDA Tile C++ 需要 C++20 与 -enable-tile；支持范围以当前 NVCC 文档为准。
• 13.4 文档说明 Tile code generation 面向 sm_80 及以上，但特定 builtin 可能要求更高架构。
• 分离编译对 Tile 代码仍有限制；不要假设与普通 SIMT RDC 完全一致。
• NVRTC 编译 Tile 需要显式选项，并通过 Tile IR/相应模块加载路径处理。
• 生成的 fatbinary 可能同时含 SIMT 与 Tile 部分，部署时需测试目标 Driver 的兼容性。
• API 快速演进，缓存 key 和构建产物必须包含 Toolkit/编译器版本。

决策树

要写自定义 GPU kernel？
├─ 已有 CUDA-X/框架算子满足 → 调库
└─ 不满足
   ├─ 规则 dense tensor / matmul / reduction？
   │  ├─ 是 → 评估 CUDA Tile；对照库和 SIMT
   │  └─ 否 → SIMT/C++ 或 Python SIMT DSL
   ├─ 需要逐线程通信、复杂原子或不规则队列？ → SIMT
   └─ 需要同一程序混用？ → host 端串联 Tile 与 SIMT kernel

练习

1. 用 SIMT 和 Tile C++ 各实现向量加法，覆盖 n=0、小于一个 tile 和非整除长度。
2. 实现二维 bias + ReLU，并改变 tile shape，记录编译时间、kernel 时间和边界浪费。
3. 写稳定 softmax tile kernel，与框架实现比较误差、吞吐和不同列数下的性能。
4. 用 profiler 验证一个 Tile matmul 是否采用目标矩阵指令和期望的数据移动路径。

官方参考

• CUDA Programming Guide: Writing Tile Kernels
• CUDA Tile C++ API Reference
• NVCC: Tile Compilation in CUDA
• NVRTC: Tile Programming
• CUDA Tile Documentation