CUDA Tile 优化提示详解：C++ 与 Python 双语言实战

一、优化提示

Optimization Hints，即“优化提示”，是 CUDA Tile 的一种元数据。它可以附加到源代码结构中，包括 Tile 的内核函数、加载及存储调用点等（前面文章都已分析过），目的是指导编译器的代码生成。理解起来也很简单：语言中的元数据多与编译器高度相关，本质上是编译器下放到代码中的“探针”。

这类优化提示对程序本身是透明的——它不会改变内核编译和运行的结果，可以放心地自由增加、减少或更改，不会对代码功能产生影响。编译器也可能直接忽略这些提示，不做任何反应，这有点类似 C++ 中的 inline 关键字。

优化提示的意义在于，它只作用于编译过程的优化阶段，不改变语义。它是一种建议而非强制要求，并且受到硬件平台的限制。如果出现冲突，编译器会自动选择最优策略。

二、属性

对于这类优化提示，存在两个通用属性：

1. 提示是独立的单个构造
   这意味着提示只应用于它附加到的特定代码块（如内核函数或调用表达式），无法影响其他代码。

2. 提示可以根据架构指定
   提示与硬件强相关。不同的 GPU 设备可以使用不同的值，当然也可以统一赋值为相同的值。

需要注意的是，cutile::hint 是 NVIDIA CUDA Tile 库特有的扩展，需要确保环境中的编译器和 CUDA 版本兼容。

CUDA Tile 面向两类编程语言时，使用方式有所不同：C++ 通过最新的属性（Attribute）来设置，可以将其安排到相关声明或语句上；Python 则使用内核装饰器的关键字参数和单个内存操作调用点的关键字参数进行操作。

三、C++ 中的属性

在 CUDA 13.3 及以上版本中，Tile 提供了扩展的 cutile::hint 属性，用于施加优化提示。基本表示如下：

[[ cutile::hint(arch, kind1=value1, kind2=value2, ...) ]]

第一个参数是目标架构，编码为与 CUDA_ARCH 宏相同的整数（例如 900 对应 sm_90，1000 对应 sm_100）。特殊值 0 表示与架构无关的提示，适用于所有目标架构。其余参数是 kind=value 对，指定提示种类及其值。

使用 cutile::hint 属性时，必须将其放在应用目标之前：

• 内核函数：将属性放在函数声明上

// 提示：每个 CGA（Cooperative Group Array）包含 4 个 CTA（Thread Blocks）
[[ cutile::hint(0, num_cta_per_cga = 4) ]]
__global__ void my_kernel(float* data, int n)
{
    // 内核逻辑
}

• 内存操作：如 ct::load、ct::store 和 ct::partition_view 的加载/存储，属性放在包含调用的表达式语句上

// 提示：在此处禁用 TMA，使用传统加载指令
[[ cutile::hint(0, allow_tma = false) ]]
auto tile_data = ct::load<float>(global_ptr, tile_shape);

• 循环结构：将属性放在 for、while 或 do-while 关键字之前

// 提示：建议编译器完全展开此循环
[[ cutile::hint(0, unroll = true) ]]
for (int i = 0; i < TILE_SIZE; ++i) {
    // 循环体
}

• 变量或对象：将属性放在变量或对象声明之前

// 提示：将此共享内存数组对齐到 128 字节边界
[[ cutile::hint(0, alignment = 128) ]]
__shared__ float shared_buffer[256];

更多细节可以参阅 CUDA 官方相关文档。

四、Python 中的装饰器参数和调用点关键字

Python 与 C++ 不同，它通过两种方式给出提示：

1. 内核级提示
   作为 @ct.kernel(…) 装饰器的关键字参数传入。编译后的内核对象还提供 .replace_hints( **hints) 方法，返回具有覆盖提示的新内核。新内核拥有自己的 JIT 缓存，这使得 replace_hints 成为自动调优循环的自然构建块。

from cuda.tile import ByTarget

# 示例：为不同架构设置不同的 num_ctas
@ct.kernel(num_ctas=ByTarget(sm_100=8, sm_120=4, default=2))
def my_kernel(x):
    ...

2. 单次调用提示
   作为内存操作调用点的关键字参数，适用于 ct.load / ct.store、TiledView.load / TiledView.store，以及 ct.gather / ct.scatter。

# 伪代码示例，具体参数需参考最新文档
a_tile = ct.load(A, index=(bidx, k), shape=(tm, tk), cache_hint="global")
ct.store(C, index=(bidx, bidy), tile=accumulator, sync_policy="async")

对于平台架构的值，可使用 cuda.tile.ByTarget(*, default=…, sm_XXX=…, sm_YYY=…) 来打包。架构键必须为 "sm_" 形式的字符串（例如 "sm_100" 或 "sm_120"）。普通（非 ByTarget）值适用于所有目标，等同于 C++ 中 arch=0 的架构无关提示。

五、种类

CUDA Tile 中的优化提示分为以下几类：

1. 每簇 CTA 数量
   C++ 使用 num_cta_in_cga（内核属性），Python 使用 num_ctas（@ct.kernel 装饰器参数）进行配置。允许设置的值为 1、2、4、8、16；在 sm_80 上只有 1 适用。
   它表示编译器启动内核时，每个协作组数组（CGA）中应优先使用的协作线程数组（CTA）数量。

2. 占用率
   C++ 使用 occupancy（内核属性），Python 使用 occupancy（@ct.kernel 装饰器参数）配置。允许设置的值为闭区间 [1, 32] 内的任意整数。
   它表示每个流多处理器（SM）上的目标活跃 CTA 数量。编译器将其视为建议，并会在代码生成时尽量遵守。

3. 内存访问延迟
   C++ 使用 latency（包含调用的表达式语句上的属性），Python 使用 latency（调用点的关键字参数）配置。它应用于 Tile 空间加载和存储（C++ 中的 ct::partition_view；Python 中的 Array.tiled_view 和 ct.load / ct.store）以及 gather / scatter（C++ 中使用指针 Tile 的 ct::load / ct::store；Python 中的 ct.gather / ct.scatter）。
   可设置的值为闭区间 [1, 10] 内的任意整数，其中 1 表示轻量 DRAM 流量，10 表示重度流量。较大的值通常使编译器调度更大的预取深度。

4. TMA 设置
   C++ 使用 allow_tma（包含调用的表达式语句上的属性），Python 使用 allow_tma（调用点的关键字参数）配置。它仅用于 Tile 空间加载和存储（C++ 中的 ct::partition_view；Python 中的 Array.tiled_view 和 ct.load / ct.store）。注意，gather 和 scatter 操作不接受此优化提示。
   可设置的值为 true / false（C++）或 True / False（Python）。默认允许使用 TMA；将此提示设为 false / False 可指示编译器在支持 TMA 的硬件上不要将此特定加载或存储降级为 TMA。

需要说明的是，这些提示在 C++ 和 Python 之间是共享的。每个提示在两种语言中的名称只是同一底层提示的不同拼写；其他方面（应用位置、取值及含义）则完全相同。

六、例程

下面看一个完整例程：

import argparse
import cupy as cp
import cuda.tile as ct
from cuda.tile import ByTarget

# 带内核级提示的向量拷贝
#   - num_ctas : 每簇 CTA 数量，按架构分别指定
#   - occupancy: 每 SM 目标活跃 CTA 数
@ct.kernel(
    num_ctas=ByTarget(default=1, sm_90=4, sm_100=8),   # 按架构指定
    occupancy=8,                                        # 所有架构统一值
)
def copy_with_kernel_hints(in_, out, TILE: ct.Constant[int]):
    bid = ct.bid(0)

    # 单次调用提示：此加载为带宽密集型，建议更大的预取深度
    tile = ct.load(in_, index=(bid,), shape=(TILE,), latency=8)

    # 单次调用提示：禁止将此存储降级为 TMA
    ct.store(out, index=(bid,), tile=tile, allow_tma=False)

# 无提示的基准版本（用于对比 / 自动调优起点）
@ct.kernel
def copy_baseline(in_, out, TILE: ct.Constant[int]):
    bid = ct.bid(0)
    tile = ct.load(in_, index=(bid,), shape=(TILE,))
    ct.store(out, index=(bid,), tile=tile)

# 多种提示
@ct.kernel(
    num_ctas=ByTarget(sm_90=2, sm_100=4, sm_120=4),
    occupancy=ByTarget(default=4, sm_100=8),
)
def copy_all_hints(in_, out, TILE: ct.Constant[int]):
    bid = ct.bid(0)

    # 轻量 DRAM 流量
    tile = ct.load(in_, index=(bid,), shape=(TILE,), latency=1)

    # 允许 TMA（默认行为，此处显式写出以作演示）
    ct.store(out, index=(bid,), tile=tile, allow_tma=True)

# 辅助函数
def launch_copy(kernel, in_, out, tile_size, num_ctas=None):
    """在给定内核上启动向量拷贝，可选地在启动前覆盖 num_ctas 提示。"""
    if num_ctas is not None:
        kernel = kernel.replace_hints(num_ctas=num_ctas)

    stream = cp.cuda.get_current_stream()
    grid = ((in_.size + tile_size - 1) // tile_size,)
    ct.launch(stream, grid, kernel, (in_, out, tile_size))
    stream.synchronize()
    return out

def main():
    parser = argparse.ArgumentParser(
        description="演示 CUDA Tile 优化提示与运行时自动调优。"
    )
    parser.add_argument("--size", type=int, default=2**16,
                        help="向量元素总数")
    parser.add_argument("--tile", type=int, default=256,
                        help="Tile 大小（2 的幂）")
    args = parser.parse_args()

    if args.size <= 0:
        raise SystemExit("--size 必须为正整数")

    stream = cp.cuda.get_current_stream()
    rng = cp.random.default_rng(0)
    x = rng.random(args.size, dtype=cp.float32)

    # 1. 使用带内核级提示的内核
    y1 = cp.empty_like(x)
    launch_copy(copy_with_kernel_hints, x, y1, args.tile)
    cp.testing.assert_array_equal(y1, x)
    print("copy_with_kernel_hints 通过（num_ctas 按架构、occupancy=8）")

    # 2. 使用基准内核
    y2 = cp.empty_like(x)
    launch_copy(copy_baseline, x, y2, args.tile)
    cp.testing.assert_array_equal(y2, x)
    print("copy_baseline 通过（无提示）")

    # 3. 使用综合提示内核
    y3 = cp.empty_like(x)
    launch_copy(copy_all_hints, x, y3, args.tile)
    cp.testing.assert_array_equal(y3, x)
    print("copy_all_hints 通过（latency=1、allow_tma=True）")

    # 4. 运行时自动调优：在基准内核上覆盖 num_ctas
    #    replace_hints 返回新内核，拥有独立的 JIT 缓存，
    #    不会影响原始 copy_baseline 的缓存。
    y4 = cp.empty_like(x)
    launch_copy(copy_baseline, x, y4, args.tile, num_ctas=4)
    cp.testing.assert_array_equal(y4, x)
    print("copy_baseline.replace_hints(num_ctas=4) 通过")

    # 5. 简单自动调优循环：测试多个 num_ctas 值
    results = []
    for num_ctas in (1, 2, 4, 8):
        y_try = cp.empty_like(x)
        tuned = copy_baseline.replace_hints(num_ctas=num_ctas)
        launch_copy(tuned, x, y_try, args.tile)
        cp.testing.assert_array_equal(y_try, x)
        results.append(num_ctas)
        print(f"   autotune 测试 num_ctas={num_ctas}: 正确")

    print(f"\n所有 {len(results)} 种 num_ctas 配置均验证通过。")

if __name__ == "__main__":
    main()

七、总结

不难推测，由于当前运行环境较为复杂，优化提示只是给编译器根据不同情况提供优化指示，而无法强制编译器执行指定优化。这也符合实际开发场景——不能因为一个局部指定就妨碍整体的优化策略。