Tile Kernel 結構與啟動 (Tile Kernel Structure)

重點總覽

CUDA Tile 提供有別於 SIMT 的 kernel 撰寫方式：programmer 以整個 block（tile）為單位思考（load 一整塊 tile、對整塊做運算、再 store），把最底層的 thread 平行對應交給 compiler。Kernel 外圍的 host 程式（cudaMalloc、cudaMemcpy、launch 排序）與 SIMT 完全相同，只有 kernel 內部寫法改變。Python 透過 cuda.tile 套件，C++ 從 CUDA Toolkit 13.3 起提供，兩者共用 compiler 後端 CUDA Tile IR，慣例上 API 別名為 ct。

SIMT 思維                          Tile 思維
┌─────────────────────┐           ┌─────────────────────┐
│ idx = blockIdx*N +  │           │ load 一整塊 tile     │
│       threadIdx     │           │   (多個元素)         │
│ load 我這一格元素   │   vs.     │ 對整塊 tile 做運算   │
│ 對單一元素運算      │           │ store 整塊 tile      │
│ store 我這一格      │           │ (thread 對應 → 編譯器)│
└─────────────────────┘           └─────────────────────┘
   programmer 管 thread              compiler 管 thread

API 別名與引入

慣例上 tile API 在兩種語言皆別名為 ct，後續所有 ct. / ct:: 前綴都指 tile API。

#include "cuda_tile.h"
namespace ct = cuda::tiles;   // C++：API 位於 cuda::tiles namespace

import cuda.tile as ct        # Python：API 位於 cuda.tile 模組

• C++ 的 tile API 由 cuda_tile.h header 暴露，位於 cuda::tiles namespace。
• Python 的 tile API 位於 cuda.tile 模組。
• 兩語言共用 compiler 後端 CUDA Tile IR，因此執行語意相同。
• Tile 程式相較 SIMT 提供更簡單的途徑去使用近期 NVIDIA GPU 的效能功能，如 TMA（tensor memory accelerator） 單元與 tensor cores。

2.4.1 Kernel 與 Function 宣告

Tile kernel 是 GPU entry point，每個 block 執行一次；tile function 可被 tile kernel 或另一個 tile function 呼叫，但本身不是 entry point。與 SIMT kernel 一樣，tile kernel 不能從 host code 直接呼叫，必須 launch。

#include "cuda_tile.h"
// Tile kernel entry point：不能直接呼叫，必須 launch
__tile_global__ void my_kernel(float* a, float* b, float* c) { ... }

// Tile function：可被 tile kernel 與 tile function 呼叫
__tile__ float helper(float x, float y) { return x + y; }

import cuda.tile as ct
@ct.kernel                      # entry point
def my_kernel(a, b, c): ...

@ct.function                    # 可省略：被 tile code 呼叫的函式會自動編成 tile code
def helper(x, y):
    return x + y

• Array 與 scalar 參數的傳遞方式與 SIMT kernel 相同。
• Tile code 與 SIMT code 可共存：單一 .cu 檔可同時定義 __tile_global__ 與 __global__ kernel，單一 host 程式可同時 launch 兩者。
• Python 中 @ct.function 是可省略的：任何被 kernel 呼叫的函式都會自動編譯為 tile code。
• Python 的 array 參數接受任何暴露 DLPack 或 CUDA Array Interface 的 device-resident array（例如 PyTorch tensor、CuPy array）；scalar 參數直接傳入。

2.4.2 啟動 Kernel

Tile kernel 啟動於一個 tile blocks 的 grid，如同 SIMT kernel 啟動於 thread blocks 的 grid。Programmer 指定 grid 形狀（最多三維）。從 programmer 視角，每個 tile block 由「單一邏輯 thread」執行；block 內的平行由 compiler 管理。

// C++：沿用 SIMT 的 triple-chevron；第二個 chevron 參數必須是 1
my_kernel<<<dim3(num_blocks_x, num_blocks_y), 1>>>(a, b, c);
//                                            ↑ 第二參數（per-block thread 數）必須是 1

import torch
stream = torch.cuda.current_stream()          # CUDA stream 物件
grid = (num_blocks_x, num_blocks_y, 1)         # tile-block grid (x, y, z)
ct.launch(stream, grid, my_kernel, (a, b, c))  # 四個位置參數

• C++：第一個 chevron 參數是 grid 形狀（tile block 數量）；第二個原本是 SIMT 的 per-block thread 數，但 tile kernel 由 compiler 內部決定 thread 數，故此參數必須為 1。
• Tile kernel 同時也是普通 CUDA kernel，可透過 runtime 既有的 cudaLaunchKernel 與 cudaLaunchKernelEx 以相同的 grid, 1 設定啟動——便於整合進已用這些 API 驅動 launch 的程式庫。
• Python：ct.launch 取四個位置參數——CUDA stream、grid tuple（各維 tile block 數）、kernel 物件、kernel 參數 tuple。

2.4.2.1 Grid-Sizing Pattern（只指定 grid，thread 數交給 compiler）

常見模式是啟動足夠多的 block 以覆蓋整個 array，其中最後一個 block 可能在一或多維上超出 array 大小。用 ceil（向上取整）除法計算 block 數。

int num_blocks = (N + tile_size - 1) / tile_size;  // ceil 除法 → 覆蓋不滿的尾端
kernel<<<num_blocks, 1>>>(in, out, N);

import math
grid = (math.ceil(N / TILE),)                      # ceil 除法 → 覆蓋不滿的尾端
ct.launch(stream, grid, my_kernel, (arr_in, arr_out, TILE))

N = 10, TILE = 4  →  num_blocks = ceil(10/4) = 3
 array:  [0 1 2 3][4 5 6 7][8 9 . .]
 block:   block0   block1   block2  ← 最後一個 block 超出 array（尾端 2 格無效）

2.4.3 查詢 Block 位置

每個 block 需知道自己在 grid 中的位置，才能決定要處理 array 的哪一部分。SIMT 中 programmer 結合 blockIdx 與 threadIdx 算出 global thread index；tile code 只需要 block index，所有 block 內的 thread-level indexing 由 compiler 處理。

__tile_global__ void my_kernel(float* a, float* b, float* c) {
    namespace ct = cuda::tiles;
    int bid_x = ct::bid().x;            // .x 方向的 block index
    int bid_y = ct::bid().y;            // .y 方向的 block index
    int num_x = ct::num_blocks().x;     // .x 方向 block 總數
}

@ct.kernel
def my_kernel(a, b, c):
    bid_x = ct.bid(0)         # axis 0 的 block index
    bid_y = ct.bid(1)         # axis 1 的 block index
    num_x = ct.num_blocks(0)  # axis 0 的 block 總數

• C++ ct::bid() 回傳含三維 block index 的 uint3；ct::num_blocks() 回傳含各維 block 總數的 dim3（依 launch 參數決定）。
• Python ct.bid(axis) 取 axis（0、1、2）回傳該軸 block index（int32 scalar）；ct.num_blocks(axis) 回傳該軸 block 總數——適合用於 bounds check 與 loop 計數。

2.4.4 建立 Tile

Tile 是 tile kernel 真正操作的對象：一個固定大小、多維的 scalar 元素陣列，其 shape 與 element type 在編譯期已知。

• 每一維必須是 2 的次方。
• Tile 具 value semantics：複製 tile 會複製其元素，兩份副本完全獨立；但因 compiler 控制 tile 在 hardware 內的表示方式，複製成本低。
• Programmer 不需也不能 為 tile 配置或釋放記憶體。
• 建立方式：(1) 從 array load 資料；(2) 用 factory function 產生填滿特定 pattern 的 tile。

C++ tile 型別是顯式的：ct::tile<T, ct::shape<dims...>>，T 是元素型別，ct::shape<dims...> 把各維大小編碼成 template 引數（編譯期已知）。

__tile__ void factories() {
    namespace ct = cuda::tiles;
    using i32x8   = ct::tile<int,   ct::shape<8>>;     // 1-D：8 個 int
    using f32x4x4 = ct::tile<float, ct::shape<4, 4>>;  // 2-D：4x4 float

    auto z      = ct::zeros<f32x4x4>();    // 全 0
    auto o      = ct::ones<f32x4x4>();     // 全 1
    auto filled = ct::full<f32x4x4>(3.14f);// 全 3.14
    auto seq    = ct::iota<i32x8>();       // {0,1,2,3,4,5,6,7}
}

@ct.function
def factories():
    zeros  = ct.zeros((64, 64), dtype=ct.float32)   # 64x64 全 0.0
    ones   = ct.ones((128,),   dtype=ct.float16)    # 128 元素全 1.0
    filled = ct.full((32, 32), 3.14, dtype=ct.float32)  # 32x32 全 3.14
    seq    = ct.arange(8, dtype=ct.int32)           # [0,1,2,3,4,5,6,7]

• C++ 因 shape 是型別的一部分，永遠在編譯期已知；範例慣用 using alias（如 using f32x4x4 = ct::tile<float, ct::shape<4, 4>>）讓 call site 可讀。
• Python 中 shape tuple 與 dtype 都是編譯期值；Python literal（如 (64, 64)、ct.float32）自然滿足，也可用 Constant-annotated 參數提供（見下節）。產生的 tile 暴露 .shape、.dtype、.ndim 屬性，反映其編譯期屬性。

2.4.5 Compile-Time Constants（編譯期常數）

Tile compiler 會為每一種 tile shape、data type 與其他結構參數的組合產生專屬機器碼。因此任何影響產生碼的值都必須在編譯期已知——也就是 tile 的 shape 與 dtype 必須是編譯期值。除了直接寫 literal，shape 也能透過 kernel 介面以「編譯期已知值」傳入。

2.4.5.1 Python Constant[T]

在 kernel 參數上加 ct.Constant[T] type hint，將其標記為 constant-embedded：該參數在 kernel 內的每一次使用，都如同把 literal 值寫在該處一樣（compiler 看到的是具體 literal，並據此產生 specialized code）。

@ct.kernel
def my_kernel(TILE: ct.Constant[int]):
    # TILE 被 constant-embedded：compiler 看到其 literal 值（如 128）並產生專屬碼，
    # 此處 TILE 驅動 factory tile 的 shape。
    zeros = ct.zeros((TILE,), dtype=ct.float32)

• 型別引數可省略：ct.Constant（不帶型別）可嵌入任意型別的常數。
• 最常用於整數 ct.Constant[int]，用在驅動 tile shape 與 loop bound 的參數上。

2.4.5.2 C++ integral_constant 與 _ic Literal

CUDA Tile C++ 透過 ct::integral_constant 表達編譯期值——其數值編碼在型別本身。ct::literals namespace 的 _ic literal 提供簡寫：0_ic 產生一個 ct::integral_constant<0> 值。

接受編譯期值的 API 同時接受兩種形式：non-type template parameter（NTTP）形式，與 _ic literal 形式。

__tile__ void concat_demo() {
    namespace ct = cuda::tiles;
    using namespace ct::literals;
    using T = ct::tile<int, ct::shape<4, 8>>;
    T lhs = ct::full<T>(0);
    T rhs = ct::full<T>(1);
    auto a = ct::cat<0>(lhs, rhs);       // NTTP 形式（軸 0 是編譯期值）
    auto b = ct::cat(lhs, rhs, 0_ic);    // _ic 形式（同一個編譯期軸）
}

ct::extents 與 ct::shape 各有 NTTP 形式與 brace 形式。與 NTTP 形式不同，brace 形式接受 runtime 值——當一或多個維度只在 launch 時才知道時，就用它：編譯期維度寫 _ic literal、runtime 維度寫普通變數。ct::tensor_span、ct::partition_view（見 Tile-Space Loads and Stores）即用此形式包裝這類 array。

auto shape2d = ct::extents{8_ic, length};  // 8 是編譯期；length 是 runtime

編譯期值的兩種寫法（C++）
 NTTP 形式：  ct::cat<0>(...)            ct::extents<std::uint32_t, 4, 8>
 _ic  形式：  ct::cat(..., 0_ic)         ct::extents{8_ic, length}
                                                    ↑      ↑
                                              編譯期   runtime（brace 形式才行）

考試/測驗重點

Related Notes

• 02-Programming-GPUs/06-SIMT-Basics-and-Thread-Hierarchy
• 02-Programming-GPUs/01-CUDA-Cpp-Kernels-and-Launch
• 02-Programming-GPUs/11-Tile-Load-Store-and-Control-Flow
• 02-Programming-GPUs/12-Tile-Operations-and-Primitives
• 02-Programming-GPUs/13-Tile-Atomics-and-Optimization
• 02-Programming-GPUs/14-Async-Streams-and-Events
• 01-Introduction-to-CUDA/03-Tile-Programming
• 02-Programming-GPUs/Practice-Programming-GPUs
• 00-Dashboard/Exam-Traps