昨天介紹了從工作組織、硬體執行、資料存取三個角度拆解 NVIDIA GPU 的 grid、block、warp、thread 是工作，SM 是執行工作的硬體，register、shared memory、global memory 則是資料位置。

今天我會先用向量加法看 CUDA programming model，再看 cuBLAS、cuDNN、CUTLASS 替開發者處理哪些工作，接著用同一個 tile-based 視角比較 CuTe、cuTile 與 Triton 這些 domain-specific languages（DSLs），之後重點會放在 Triton compile flow。

本篇大綱

• 先看 CUDA 的 host／device、kernel、thread hierarchy 與 memory hierarchy
• 簡單介紹 cuBLAS、cuDNN 與 CUTLASS
• 接著比較 CuTe DSL 與 cuTile 兩種 Python DSL
• 最後介紹 Triton，比較 CUDA、cuTile 與 Triton 的程式模型

CUDA programming model：先理解 GPU 如何執行 kernel

CUDA（Compute Unified Device Architecture）是 NVIDIA 在 2006 年推出的 GPU 運算平台與 programming model。CUDA 出現以前，GPU 主要藉由圖形 API 使用，開發者若想把一般運算放到 GPU，必須把問題轉成圖形 pipeline 能處理的形式， CUDA 提供一套通用運算介面，開發者可以直接撰寫 GPU kernel，再由 CPU 端配置記憶體與啟動工作。

CUDA 涵蓋的範圍比程式語言更廣。平常所說的 CUDA C++ 包含 C++ language extensions、runtime API、driver API 與編譯工具鏈。__global__、threadIdx 和 <<<...>>> 屬於程式介面，下列名詞會在後續的 compiler flow 一直出現。

nvcc 可以保留 PTX，也可以呼叫 ptxas 產生 cubin。PTX 提供跨部分 GPU 世代的 virtual ISA。
cubin 已經針對 sm_XX 目標組譯，距離 GPU 最終執行的 machine code 更近。

CUDA C++ source
  -> NVIDIA CUDA Compiler (nvcc)
      -> Parallel Thread Execution (PTX)
          -> PTX assembler (ptxas)
              -> CUDA binary (cubin)
                  -> CUDA Driver
                      -> GPU 執行 kernel

這個生態的重要性在於，上層的 cuBLAS、cuDNN、CUTLASS 和 Python DSL 雖然提供不同抽象，最後還是要建立 GPU 可執行程式、由 CUDA Driver 啟動，並受到 SM、warp 與 memory hierarchy 限制。

Host、device 和 kernel launch

CUDA 把系統看成 host 與 device，Host 通常是 CPU，負責配置記憶體、準備資料與啟動工作，device 是 GPU，負責執行大量平行運算。GPU 上執行的函式稱為 kernel。

下面是一個向量加法 kernel

__global__ void add(float* a, float* b, float* c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

Host 端會指定 grid 與 block 的大小，再啟動 kernel

int threads_per_block = 256;
int blocks = (n + threads_per_block - 1) / threads_per_block;
add<<<blocks, threads_per_block>>>(a, b, c, n);

<<<blocks, threads_per_block>>> 是 execution configuration。這次 launch 會建立一個 grid，裡面有 blocks 個 thread blocks，每個 block 有 256 個 threads。

Host code
  -> 準備 device memory
  -> launch kernel <<<grid, block>>>
      -> GPU 執行大量 threads
  -> 等待或繼續安排其他工作

Kernel launch 對 host 通常是 asynchronous。CPU 啟動 kernel 後可以繼續工作，需要讀回結果或量測時間時，才在正確的位置同步。

Grid、block、thread 與 warp

CUDA 的 thread hierarchy 可以先記成

Grid
  -> Thread Block
      -> Thread
          -> hardware 以 warp 為單位執行

每個 thread 會根據自己的位置決定要處理哪一筆資料

int i = blockIdx.x * blockDim.x + threadIdx.x;

• threadIdx.x：thread 在 block 內的位置。
• blockIdx.x：block 在 grid 內的位置。
• blockDim.x：每個 block 的 thread 數量。
• gridDim.x：grid 的 block 數量。

同一個 block 的 threads 會在同一個 SM 上執行，可以使用 shared memory 交換資料，也能用 __syncthreads() 做 block 內同步。不同 blocks 的排程順序沒有保證，因此一般 kernel 不能假設另一個 block 已經完成。

CUDA 採用 SIMT（Single Instruction, Multiple Threads）模型。開發者撰寫單一 thread 要做的工作，硬體會把 threads 組成 warp 執行。同一個 warp 走到不同分支時會發生 branch divergence，部分 lanes 必須等待其他路徑完成。

Memory hierarchy

昨天介紹的 HBM、L2、shared memory 與 registers，到了 CUDA 會變成可以在程式中安排的資料位置。這裡先保留寫 kernel 直接會碰到的三種，完整的可見範圍、生命週期與實體位置可以回到昨天的 memory hierarchy 表格複習

Global memory 容量大，但存取延遲較高。相鄰 threads 若能存取相鄰位址，硬體較容易形成 coalesced memory access。Shared memory 位在 SM 上，適合暫存 GEMM tile；使用時還要處理同步、容量與 bank conflict。Registers 最接近運算單元，但使用量過高可能降低同一個 SM 能同時執行的 warps 數量。

一些常見效能問題

前面的向量加法很短，仍然已經出現 GPU kernel 常見的效能問題

• blockDim.x 會影響每個 block 有多少 threads
• i 的計算會決定 global memory access 是否連續
• 分支會影響同一個 warp 的 threads 能否一起前進
• 每個 thread 做的工作太少，kernel launch overhead 可能變明顯
• 如果是 matmul，還要自己處理 tiling、shared memory、Tensor Core、pipeline

CUDA C++ 可以直接決定 thread mapping、shared memory layout、synchronization 和 inline PTX。要把 GEMM 或 attention kernel 寫快，還要一起處理資料搬移、運算切分與硬體 pipeline。接下來介紹的函式庫與 DSL，都是在不同位置接手這些工作。

簡單認識 cuBLAS、cuDNN 與 CUTLASS

理解 CUDA programming model 後，再看 NVIDIA 提供的函式庫，會比較容易理解它們替開發者處理了哪些工作。表格中的 GEMM 是 general matrix-matrix multiplication，也就是通用矩陣乘法。

cuBLAS／cuBLASLt

BLAS（Basic Linear Algebra Subprograms）是一套已建立數十年的線性代數介面規範。Level 1 處理 vector–vector 運算，Level 2 處理 matrix–vector 運算，Level 3 處理 matrix–matrix 運算。深度學習的 linear layer 與 attention 會用到大量矩陣乘法，因此 Level 3 的 GEMM 特別重要。

cuBLAS 是 NVIDIA 在 CUDA 上的 BLAS 實作。呼叫者提供輸入指標、matrix shape、transpose 設定、資料型別與 CUDA stream，函式庫在內部選擇適合的 kernel。對 AI 工程來說，cuBLAS 最常處理的是 GEMM

C = A @ B

標準 GEMM 通常可以先用 cuBLAS 建立效能 baseline。這類 API 的優點是介面穩定，且不必自行處理 thread mapping、Tensor Core 指令與完整 GEMM pipeline。限制也很明確：你是呼叫 NVIDIA 已實作的運算，無法像撰寫 kernel 那樣任意改寫運算內部。

cuBLASLt 是以 matmul 為中心的較彈性介面。開發者可以用 descriptors 表達 matrix layout、資料型別、batched matmul 與 epilogue，再由 heuristic 找出候選演算法。例如 GEMM 後緊接 bias 或 activation 時，支援的 epilogue 可以減少額外 kernel launch 與中間資料寫回 global memory 的機會。

framework 或 application
  -> 建立 matmul descriptors
      -> heuristic 回傳候選演算法
          -> 呼叫選定的內部 kernel
              -> 送進 CUDA stream

cuDNN

cuBLAS 擅長線性代數，但深度學習模型還有 convolution、pooling、normalization、activation 與 attention 等運算。這些運算會隨 tensor layout、資料型別、shape 與硬體不同，產生許多實作組合。cuDNN（CUDA Deep Neural Network library）就是 NVIDIA 針對 deep neural network 建立的 GPU primitives 與後端函式庫。PyTorch 或 TensorFlow 這類框架可以在後端呼叫它，上層使用者通常不會直接看到裡面的 CUDA kernel。

早期的 cuDNN 使用者常以單一 operation 的 API 呼叫 convolution 或 activation。cuDNN 8.0 引入 Graph API 後，框架可以先把多個 operations 與 tensor 關係建成 operation graph，再請 heuristics 找出候選 engine configuration，最後建立 execution plan。編譯器或框架因此能提供更大的運算範圍，讓函式庫判斷是否能使用 fused implementation。

operations + tensors
  -> operation graph
      -> heuristics 選擇 engine configuration
          -> execution plan
              -> 在 CUDA stream 執行

以 attention 為例，運算包含 matmul、scaling、mask、softmax 與資料搬移。框架若能把這段 graph 交給 cuDNN，就有機會使用 NVIDIA 已針對支援硬體與資料型別調校的實作。如果 operation 形狀非常特殊，或者需要研究新的資料流與量化方法，開發者才會繼續往 CUTLASS 或自訂 kernel 移動。

CUTLASS C++：用 CUDA C++ templates 組合 kernel

CUTLASS（CUDA Templates for Linear Algebra Subroutines and Solvers）從 2017 年起提供 CUDA C++ template abstractions，起點是把 NVIDIA 內部高效 GEMM 實作中重複出現的結構變成可重用元件。直接呼叫 cuBLAS 時，開發者選擇一個已建立的 operation；使用 CUTLASS 時，開發者可以調整這個 operation 內部的 tile shape、memory layout、multiply-accumulate（MMA）、pipeline 與 epilogue。

C++ templates 適合這類工作，因為 data type、tile shape、layout 與 GPU architecture 可以在編譯時期成為專用實作。不需要的分支可以在編譯時期移除，也能選擇特定硬體支援的 Tensor Core 指令。因此，CUTLASS 是一套產生與組合 CUDA kernel 的 C++ template library，提供的粒度比單一 GEMM 函式更細。

它把高效 GEMM 常見的 tile、資料搬移、MMA、pipeline 與 epilogue 拆成可組合元件。開發者選擇 shape、layout、資料型別與硬體架構後，C++ compiler 會藉由 template instantiation 產生對應的 CUDA kernel。

選擇 data type、tile shape、layout、MMA、pipeline、epilogue
  ↓ C++ template instantiation
產生特定 GPU 架構使用的 CUDA kernel

CUTLASS C++ 位於直接呼叫函式庫與自行撰寫 CUDA kernel 之間。開發者仍然使用 CUDA C++，但不必自己實作所有 GEMM 元件。資料路徑可以簡化成：

global memory
  -> shared memory
      -> registers
          -> Tensor Core MMA

開發者仍要理解 CUDA 的 thread hierarchy 與 memory hierarchy，但可以重用 CUTLASS 已整理好的 building blocks，組出特定 shape、layout、資料型別或 fused epilogue 的 kernel。

代價是 C++ template 類型很長、錯誤訊息難讀，而且大量 template instantiation 會增加編譯與迭代時間。CUTLASS 3.0 把 CuTe 納入底層抽象，CUTLASS 4 再加入 Python DSL，就是為了保留這些布局與硬體元件，同時縮短開發迴圈。

CuTe DSL：用 Python 操作 CUTLASS 的底層抽象

CUTLASS 3.0 引入的 CuTe C++ 以 layout algebra 為核心，用 shape 與 stride 描述邏輯座標如何對應到記憶體位置，也可用來表達資料如何映射給 threads。這讓 tensor layout、tiled copy 與 tiled MMA 能使用同一組組合規則描述。

這種寫法提供很細的控制，卻也把 C++ template metaprogramming 的複雜度帶給 kernel 開發者。NVIDIA 於 2025 年 6 月 釋出 CUTLASS 4.0，其中首次加入 CuTe DSL（CuTe Domain-Specific Language）。

CuTe DSL 讓開發者在 Python 環境中表達 CuTe 的 layout、tensor、copy、MMA 與 pipeline。Python 負責 metaprogramming 與 just-in-time（JIT）特化，底層仍保留能對應 NVIDIA GPU 記憶體與 Tensor Core 的抽象。

CuTe DSL 還是會讓你面對 thread 與資料的硬體階層。開發者需要安排 tile、memory layout、copy atom、MMA atom 與 pipeline，只是把複雜的 C++ template metaprogramming 換成 Python。它適合需要細部硬體控制，又想縮短 kernel 開發與編譯迭代時間的工作。

cuTile：讓編譯器處理 thread-level mapping

CUDA 長期以 SIMT 與 thread hierarchy 為中心，這種模型能完整暴露 GPU，也要求開發者把運算切到 threads、處理合作與同步。新 GPU 不斷加入更複雜的非同步資料搬移與 matrix instructions，同一個 tile operation 在不同架構上的理想 thread mapping 也可能改變。CUDA Tile 因此改用 tile-based programming model，開發者先描述整塊資料的運算，編譯器再負責將它映射到 block 內的 threads。

NVIDIA 在 2025 年 12 月發布的 CUDA Toolkit 13.1 首次推出 CUDA Tile，內容包含 CUDA Tile IR 與 cuTile。CUDA Tile IR 是 tile-based virtual ISA，cuTile 是建立在它上面的 Python DSL，使用 cuda.tile API。初始版本以 Blackwell GPU 為目標，後續 CUDA 13.x 才逐步擴大架構支援。

cuTile 程式以 logical grid of blocks 執行，每個 block 對 tile 做 load、store、矩陣乘法、reduction 與 shape transformation。開發者仍然決定 grid 如何切分整體問題，但只描述 block-level tile 運算，不直接指定 block 內每個 thread 的工作。

把前面的向量加法改用 cuTile，可以寫成下面這樣。這裡為了專注在程式模型，先假設向量長度是 TILE_SIZE 的整數倍

import cuda.tile as ct

TILE_SIZE = 256

@ct.kernel
def add_kernel(a, b, output):
    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(output, index=(block_id,), tile=a_tile + b_tile)

這裡的 block_id 選擇資料 tile，ct.load 載入整塊資料。cuTile 的 block 是 logical execution unit，tile 是資料單位；一個 block 可以操作多個 tiles。Block 內的 array operations 由 threads 共同執行，語言不暴露個別 thread，也不允許開發者在 block 內寫明確同步。

cuTile 還區分 global array 與 tile value。Global array 對應有 shape、stride 與可變儲存空間的記憶體，tile 則是 block 內的 immutable value，shape 在編譯時期已知。開發者不必先宣告 tile 一定放在 registers 或 shared memory，這些物理儲存與 thread mapping 由編譯器決定。這是 cuTile 能簡化程式的來源，也是效能分析時需要檢查 compiler output 的原因。

CuTe DSL 和 cuTile 都用 Python 寫 GPU kernel，但控制界線不同

cuTile 和 Triton 都讓開發者先描述一個 program instance 要處理的資料 tile，再由編譯器映射到底層 GPU。兩者使用不同的 API、IR 與編譯器。接著用 Triton 的向量加法對照。

Triton：用 blocked program 描述 GPU kernel

已有的 GPU 函式庫可以快速處理標準 GEMM 與常見 DNN operations，但框架開發者還是會遇到 fused operator、特殊 data layout、sparse computation 或新演算法。如果每次都改寫 CUDA C++，開發者必須同時處理運算邏輯、thread scheduling、memory coalescing 與硬體指令。Triton 的出發點是讓程式設計者描述 blocked algorithm，再由 compiler 從 block-level data flow 分析資料局部性與平行性。

Triton 是怎麼開始的？

Triton 的起點可以從原始開發者 Philippe Tillet，他有來台灣讀過書，在 2014 年取得交大碩士學位。
之後在哈佛大學研究 GPU 上的 blocked algorithms 和 compiler。有個傳聞是他在 2018 年啟動 Triton，起因是用 CUDA 撰寫矩陣乘法 auto-tuner 時遇到開發困難。2019 年，Philippe Tillet、Hsiang-Tsung Kung 與 David Cox 發表 Triton: An Intermediate Language and Compiler for Tiled Neural Network Computations，公開 Triton 的語言與 compiler 設計。他在 2020 年完成博士論文 Blocked Algorithms for Neural Networks: Design and Implementation on GPUs。

Philippe Tillet 在 2020 年加入 OpenAI，繼續全職開發 Triton compiler。OpenAI 延續改進這個系統，並於 2021 年 7 月 釋出 Triton 1.0。

Triton 同時是程式語言與 compiler

Triton 是 Python-based GPU programming language 與 compiler。這兩個身分可以分開看

• Language：它定義 @triton.jit、tl.program_id、tl.arange、tl.load、tl.store 與 tl.dot 等語法和語意，讓開發者使用 Python 語法描述 GPU kernel。
• Compiler：它讀取裝飾過的 Python function，建立 Triton IR，進行 layout、memory access、pipeline 與 target-specific lowering，最後為 NVIDIA 或其他支援的 backend 產生 GPU binary。

因此，@triton.jit function 不會像一般 Python function 那樣逐行由 interpreter 在 CPU 上處理 tensor elements。Decorator 會把 function 交給 Triton frontend，compiler 再根據資料型別、constexpr 與 meta-parameters 建立專用版本。這是 Triton 同時被稱為 language、compiler 與 JIT compiler 的原因。

程式模型主要以 blocked program 組織平行工作。CUDA 範例中的一個程式實例對應一個 thread，Triton 的一個程式實例通常一次處理一整個 data block。這個 block 是語言中可以分析的 value 集合，編譯器因此比只看單一 scalar thread 擁有更多資料存取與運算結構的資訊。

同樣是向量加法，Triton kernel 可以寫成

import triton
import triton.language as tl

@triton.jit
def add_kernel(x, y, output, n_elements, BLOCK_SIZE: tl.constexpr):
    pid = tl.program_id(axis=0)
    offsets = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
    mask = offsets < n_elements

    x_block = tl.load(x + offsets, mask=mask)
    y_block = tl.load(y + offsets, mask=mask)
    tl.store(output + offsets, x_block + y_block, mask=mask)

pid 決定這個程式實例要處理哪一段資料，tl.arange 建立一組 offsets，tl.load 與 tl.store 則一次操作一個 block of values。開發者不需要替每個 GPU thread 寫索引，編譯器會把 block-level operations 映射到 threads、warps 與硬體指令。

Triton 仍然需要 CUDA 章節建立的硬體直覺。開發者要選擇 BLOCK_SIZE、安排連續存取、減少不必要的資料搬移，並觀察 registers、shared memory、occupancy 與 Tensor Core 使用情況。抽象提高了，效能問題仍會回到同一套 GPU 硬體限制。

Triton compiler 主要接手程式實例內的 memory coalescing、shared memory management 與 thread scheduling。開發者還是要決定 grid、tiling strategy 與 program instances 之間的工作分配。

cuTile 與 Triton 的程式模型差異

cuTile 和 Triton 都用 Python 描述一塊資料要做的運算，差異可以從 block、資料模型與可調參數來看

cuTile 與 Triton 的 compile flow

cuTile：CUDA Tile IR

ct.launch() 看到實際 launch arguments 後，會為 @ct.kernel 進行 specialization，產生 CUDA Tile IR，也可表示成 TileIR bytecode。後續由 tileiras、libNVVM 與 ptxas 組成的 NVIDIA compiler toolchain 產生 cubin。

CUDA Tile IR 是 cuTile 前端與後端 toolchain 之間的主要公開邊界。開發者可以看到 tile-level program，而 tileiras 內部如何將 tile 映射到 threads、warps 與硬體指令，官方文件沒有列出和 Triton 同等細節的中間 IR 階段。

cuTile 也提供 export_kernel()。開發者可以 ahead-of-time（AOT）匯出 TileIR bytecode，或直接產生指定 GPU target 的 cubin，不一定要等到應用程式執行 ct.launch() 才開始編譯。

Triton：將 target-specific lowering 分成多個階段

Triton kernel 以 grid 呼叫時，JIT compiler 會根據 argument signature、constexpr 與 meta-parameters 建立特化版本。前端先產生 Triton IR（TTIR），這個階段還保留 tl.load、tl.store 與 tl.dot 這類 block-level operations。

進入 Triton GPU IR（TTGIR）後，target、layout、memory coalescing、num_warps、matmul acceleration 與 software pipeline 等決策開始具體化。後續再 lower 到 LLVM dialect 與 LLVM IR，產生這裡簡寫為 LLIR 的產物；NVIDIA backend 接著生成 PTX，最後由 ptxas 組譯成 cubin。

Triton 將這些階段分開後，開發者可以對照 TTIR、TTGIR、LLIR、PTX 與最後的 cubin，追蹤某個 layout 或 pipeline 決策在哪個 lowering 階段出現，我們之後會來實驗走過整個流程。

CUDA 與 Triton 比較

CUDA 和 Triton 都需要處理 tiling、memory access、data reuse 與 parallelism。主要差異在於工作切分由誰表達，CUDA 開發者從 thread 出發，Triton 開發者從 data block 出發，再由 compiler 完成較細的 thread-level mapping。

今天先走到這裡

• CUDA（Compute Unified Device Architecture）用 grid、block、thread 與 warp 組織運算；nvcc、PTX、ptxas 與 cubin 則位在編譯工具鏈的不同位置。
• cuBLAS、cuDNN 讓應用程式直接使用已最佳化的 GEMM 與 DNN primitives；CUTLASS 提供組合特殊 kernel 的 C++ building blocks。
• CuTe DSL 於 2025 年隨 CUTLASS 4.0 推出，使用 Python 暴露 layout、copy、MMA 與 pipeline；FlashAttention-4 是完整的 attention kernel 案例。
• cuTile 於 2025 年 12 月隨 CUDA 13.1 推出；它與 Triton 都從 tile／block-level operation 描述工作，但使用不同的 IR 與 compiler toolchain。
• cuTile 把 CUDA Tile IR 交給 tileiras；Triton NVIDIA backend 則依序產生 TTIR、TTGIR、LLIR、PTX 與 cubin。
• Triton 的 2019 年論文早於 OpenAI 在 2021 年開源釋出的 Triton 1.0；它同時定義 GPU programming language 與對應的 JIT compiler。

在進入 triton compile flow 介紹之前，要先學 MLIR，所以明天會先讀 MLIR Paper，理解 dialect、operation、lowering 與 pass pipeline

參考資料

• CUDA Programming Guide
  https://docs.nvidia.com/cuda/cuda-programming-guide/index.html
• cuBLAS documentation
  https://docs.nvidia.com/cuda/cublas/
• NVIDIA cuDNN documentation
  https://docs.nvidia.com/deeplearning/cudnn/
• NVIDIA CUTLASS documentation
  https://docs.nvidia.com/cutlass/latest/overview.html
• CuTe DSL overview
  https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/overview.html
• cuTile Python documentation
  https://docs.nvidia.com/cuda/cutile-python/
• Triton programming guide
  https://triton-lang.org/main/programming-guide/chapter-1/introduction.html
• Triton 2019 paper
  https://dl.acm.org/doi/10.1145/3315508.3329973
• OpenAI：Introducing Triton 1.0（2021）
  https://openai.com/index/triton/
• Triton 作者簡介：Philippe Tillet
  https://developer.nvidia.com/blog/author/philopenai/