NVIDIA CUDA 13.1 版本新增了基于 Tile 的GPU 編程模式。它是自 CUDA 發(fā)明以來 GPU 編程最核心的更新之一。借助 GPU tile kernels，可以用比 SIMT 模型更高的層級來實現(xiàn)算法。至于如何將計算任務(wù)拆分到各個線程，完全由編譯器和運行時在底層自動處理。不僅如此，tile kernels 還能夠屏蔽 Tensor Core 等專用硬件的細(xì)節(jié)，寫出的代碼還能兼容未來的 GPU 架構(gòu)。借助 NVIDIA cuTile Python，開發(fā)者可以直接用 Python 編寫 tile kernels。

什么是cuTile Python？

cuTile Python 是 CUDA Tile 編程模型在 Python 中的實現(xiàn)，基于 CUDA Tile IR 規(guī)范開發(fā)。它支持用 Python 編寫 tile kernels，以 tile-based 模型來定義 GPU kernels——既可以作為 SIMT 模型的替代，也能作為 SIMT 模型的補充。

SIMT 編程要求明確指定每個 GPU 執(zhí)行線程的任務(wù)。理論上，每個線程都能獨立運行，執(zhí)行和其他線程不同的代碼路徑。但實際應(yīng)用中，要充分發(fā)揮 GPU 性能，通常會采用單線程對不同數(shù)據(jù)執(zhí)行相同操作的算法設(shè)計思路。

SIMT 模型的優(yōu)勢是靈活性高、可定制性強，但要達(dá)到頂級性能，往往需要大量手動調(diào)優(yōu)。而 tile model 能夠幫助屏蔽部分硬件底層細(xì)節(jié)，而聚焦于更高層級的算法設(shè)計。至于 tile 算法如何拆分為線程、如何調(diào)度到 GPU 上執(zhí)行，這些工作都由 NVIDIA CUDA 編譯器和運行時自動完成。

cuTile 是專為 NVIDIA GPU 設(shè)計的并行 kernels 編程模型，核心規(guī)則有四條：

數(shù)組是最核心的數(shù)據(jù)結(jié)構(gòu)；

Tiles 是 kernels 操作的數(shù)組子集；

Kernels 是由多個 Block 并行執(zhí)行的函數(shù)；

Block 是 GPU 計算資源的子集，tiles 的操作會在各個 block 之間并行開展。

cuTile 能自動處理塊級并行、異步執(zhí)行、內(nèi)存遷移等 GPU 編程的底層細(xì)節(jié)。它可以充分利用 NVIDIA 硬件的高級特性，比如 Tensor Cores、共享內(nèi)存、Tensor 內(nèi)存加速器，而且不需要手動編寫相關(guān)代碼。更重要的是，cuTile 能跨不同 NVIDIA GPU 架構(gòu)遷移，不用重寫代碼，就能使用最新的硬件功能。

cuTile適合哪些人使用？

cuTile 面向的是需要編寫通用數(shù)據(jù)并行 GPU kernels 的開發(fā)者。目前我們重點針對 AI/ML 應(yīng)用中常見的計算類型優(yōu)化 cuTile，后續(xù)還會持續(xù)迭代——新增功能和性能特性，讓它能優(yōu)化更多類型的工作負(fù)載。

你可能會疑惑：CUDA C++ 和 CUDA Python 一直很好用，為什么還要用 cuTile 寫 kernels？關(guān)于這一點，我們在另一篇介紹 CUDA tile model 的文章里有詳細(xì)說明。簡單來說，現(xiàn)在 GPU 硬件架構(gòu)越來越復(fù)雜，我們提供這樣一層合理的抽象，就是為了能讓開發(fā)者更專注于算法本身，不用再花大量精力把算法手動適配到特定硬件上。

用 tile 模式寫代碼，既能利用 Tensor Cores 的性能，又能保證代碼兼容未來的 GPU 架構(gòu)。就像 PTX 是 SIMT 模型的底層虛擬指令集架構(gòu)（ISA），CUDA Tile IR 就是 tile-based 編程的虛擬指令集架構(gòu)。它支持用更高層級表達(dá)算法，軟件和硬件會在底層自動把這種表達(dá)映射到 Tensor Cores，助力實現(xiàn)峰值性能。

cuTilePython代碼示例

cuTile Python 代碼長什么樣？如果你學(xué)過 CUDA C++，一定接觸過經(jīng)典的向量加法 kernels。假設(shè)數(shù)據(jù)已經(jīng)從主機端拷貝到設(shè)備端，CUDA SIMT 實現(xiàn)的向量加法 kernels 如下——它接收兩個向量，逐元素相加后生成第三個向量，是最基礎(chǔ)的 CUDA kernels 之一。

__global__ void vecAdd(float* A, float* B, float* C, int vectorLength)

{

/* calculate my thread index */

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

if(workIndex < vectorLength)

{

/* perform the vector addition */

C[workIndex] = A[workIndex] + B[workIndex];

}

}在這個 kernel 里，每個線程的任務(wù)都要明確指定。而且你啟動 kernel 時，還得手動選擇要啟動的 blocks 和線程數(shù)。

再看等效的 cuTile Python 實現(xiàn)：不用指定每個線程的操作，只需把數(shù)據(jù)拆成 tiles，定義好每個 tile 的數(shù)學(xué)運算就行，剩下的工作全由 cuTile 自動處理。

cuTile Python kernel 代碼如下：

import cuda.tile as ct

@ct.kernel

def vector_add(a, b, c, tile_size: ct.Constant[int]):

# Get the 1D pid

pid = ct.bid(0)

# Load input tiles

a_tile = ct.load(a, index=(pid,) , shape=(tile_size, ) )

b_tile = ct.load(b, index=(pid,) , shape=(tile_size, ) )

# Perform elementwise addition

result = a_tile + b_tile

# Store result

ct.store(c, index=(pid, ), tile=result)

ct.bid(0) 是用于獲取（本例中）第 0 維度塊 ID 的函數(shù)，例如，它的作用相當(dāng)于 SIMT kernels 開發(fā)者使用 blockIdx.x 與 threadIdx.x。ct.load() 函數(shù)則用于從設(shè)備內(nèi)存中加載指定索引和形狀的 tile 數(shù)據(jù)，數(shù)據(jù)加載到 tiles 后，即可用于計算，所有計算完成后，ct.store() 會將 tiled 數(shù)據(jù)寫回 GPU 設(shè)備內(nèi)存。

完整實現(xiàn)代碼

接下來，我們將展示如何在 Python 中調(diào)用這個 vector_add kernels，并提供一份你可直接運行的完整 Python 腳本。以下是包含 kernels 定義與主函數(shù)的完整代碼。

"""

Example demonstrating simple vector addition.

Shows how to perform elementwise operations on vectors.

"""

from math import ceil

import cupy as cp

import numpy as np

import cuda.tile as ct

@ct.kernel

def vector_add(a, b, c, tile_size: ct.Constant[int]):

# Get the 1D pid

pid = ct.bid(0)

# Load input tiles

a_tile = ct.load(a, index=(pid,) , shape=(tile_size, ) )

b_tile = ct.load(b, index=(pid,) , shape=(tile_size, ) )

# Perform elementwise addition

result = a_tile + b_tile

# Store result

ct.store(c, index=(pid, ), tile=result)

def test():

# Create input data

vector_size = 2**12

tile_size = 2**4

grid = (ceil(vector_size / tile_size),1,1)

a = cp.random.uniform(-1, 1, vector_size)

b = cp.random.uniform(-1, 1, vector_size)

c = cp.zeros_like(a)

# Launch kernel

ct.launch(cp.cuda.get_current_stream(),

grid, # 1D grid of processors

vector_add,

(a, b, c, tile_size))

# Copy to host only to compare

a_np = cp.asnumpy(a)

b_np = cp.asnumpy(b)

c_np = cp.asnumpy(c)

# Verify results

expected = a_np + b_np

np.testing.assert_array_almost_equal(c_np, expected)

print("? vector_add_example passed!")

if __name__ == "__main__":

test()

如果已經(jīng)安裝好 cuTile Python、CuPy 等必要軟件，運行代碼很簡單，直接執(zhí)行以下命令即可：

$ python3 VectorAdd_quickstart.py

? vector_add_example passed!

恭喜！你已經(jīng)成功運行了第一個 cuTile Python 程序。

開發(fā)者工具支持

[]()cuTile kernels 的性能分析可以用 NVIDIA Nsight Compute，操作方式和 SIMT kernels 完全一樣。

$ ncu -o VecAddProfile --set detailed python3 VectorAdd_quickstart.py

生成性能分析文件后，用 Nsight Compute 圖形界面打開，按以下步驟操作：

● 選中 vector_add kernel；

● 選擇“Details”標(biāo)簽頁；

● 展開“Tile Statistics”報告板塊。

此時你會看到類似圖 1 的界面。

圖 1. Nsight Compute生成的性能分析報告，展示vector_add kernel的tile 統(tǒng)計信息

請注意，Tile Statistics 板塊包含了指定的 tile 數(shù)量、大?。ㄓ删幾g器自動選擇），以及其他 tile 專屬信息。

源碼頁面同樣支持 cuTile kernels，還能查看源碼行級別的性能指標(biāo)——這和 CUDA C kernels 的支持方式完全一致。

如何獲取 cuTile ？

運行 cuTile Python 程序，需要滿足以下環(huán)境要求：

GPU 計算能力需為 10.x 或 12.x（后續(xù) CUDA 版本會支持更多 GPU 架構(gòu)）；

NVIDIA 驅(qū)動版本需為 R580 及以上（若要使用 tile 專屬開發(fā)者工具，需升級到 R590）；

安裝 CUDA Toolkit 13.1 及以上版本；

Python 版本需為 3.10 及以上；

安裝 cuTile Python 包：執(zhí)行 pip install cuda-tile 命令即可完成安裝。

關(guān)于作者

Jonathan Bentz 領(lǐng)導(dǎo) NVIDIA 的 CUDA 技術(shù)營銷工程團隊，其團隊專注于創(chuàng)建和提供引人入勝的內(nèi)容，并與 CUDA 開發(fā)者建立聯(lián)系。Jonathan 擁有愛荷華州立大學(xué)化學(xué)博士學(xué)位和計算機科學(xué)碩士學(xué)位。

Tony Scudiero 是 CUDA 平臺的技術(shù)營銷工程師。他致力于將 CUDA 帶給各種類型和能力的開發(fā)者。在 NVIDIA 任職期間，他曾使用過大型 HPC 系統(tǒng)和應(yīng)用、實時聲學(xué)模擬 (VRWorks Audio) 和 Omniverse RTX 渲染器。