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

什么是cuTile Python？

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

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

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

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

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

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

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

Block 是 GPU 計(jì)算資源的子集，tiles 的操作會(huì)在各個(gè) block 之間并行開(kāi)展。

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

cuTile適合哪些人使用？

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

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

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

cuTilePython代碼示例

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

}

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

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

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 開(kāi)發(fā)者使用 blockIdx.x 與 threadIdx.x。ct.load() 函數(shù)則用于從設(shè)備內(nèi)存中加載指定索引和形狀的 tile 數(shù)據(jù)，數(shù)據(jù)加載到 tiles 后，即可用于計(jì)算，所有計(jì)算完成后，ct.store() 會(huì)將 tiled 數(shù)據(jù)寫(xiě)回 GPU 設(shè)備內(nèi)存。

完整實(shí)現(xiàn)代碼

接下來(lái)，我們將展示如何在 Python 中調(diào)用這個(gè) vector_add kernels，并提供一份你可直接運(yùn)行的完整 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 等必要軟件，運(yùn)行代碼很簡(jiǎn)單，直接執(zhí)行以下命令即可：

$ python3 VectorAdd_quickstart.py

? vector_add_example passed!

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

開(kāi)發(fā)者工具支持

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

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

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

● 選中 vector_add kernel；

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

● 展開(kāi)“Tile Statistics”報(bào)告板塊。

此時(shí)你會(huì)看到類似圖 1 的界面。

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

請(qǐng)注意，Tile Statistics 板塊包含了指定的 tile 數(shù)量、大小（由編譯器自動(dòng)選擇），以及其他 tile 專屬信息。

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

如何獲取 cuTile ？

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

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

NVIDIA 驅(qū)動(dòng)版本需為 R580 及以上（若要使用 tile 專屬開(kāi)發(fā)者工具，需升級(jí)到 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ù)營(yíng)銷工程團(tuán)隊(duì)，其團(tuán)隊(duì)專注于創(chuàng)建和提供引人入勝的內(nèi)容，并與 CUDA 開(kāi)發(fā)者建立聯(lián)系。Jonathan 擁有愛(ài)荷華州立大學(xué)化學(xué)博士學(xué)位和計(jì)算機(jī)科學(xué)碩士學(xué)位。

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