?
在大模型算力競(jìng)賽的下半場(chǎng)，硬件性能只是入場(chǎng)券，軟件生態(tài)的易用性是決勝的關(guān)鍵決定因素之一。昇騰（Ascend）NPU 憑借獨(dú)特的Cube/Vector（矩陣/向量）存算分離架構(gòu)，在硬件性能上表現(xiàn)強(qiáng)勁。然而，對(duì)于開(kāi)發(fā)者而言，要想完全釋放這份性能，往往面臨著“易用性”與“高性能”的艱難權(quán)衡：

使用PyTorch等上層框架開(kāi)發(fā)固然簡(jiǎn)單，自定義算子性能優(yōu)化卻難以觸達(dá)NPU特有的存儲(chǔ)層級(jí)紅利；若直接使用底層語(yǔ)言進(jìn)行編程，則需要開(kāi)發(fā)者深入理解Cube與Vector的并行關(guān)系、手動(dòng)管理復(fù)雜的多級(jí)內(nèi)存搬運(yùn)及核間同步，開(kāi)發(fā)門(mén)檻極高。

近日，TileLang AscendNPU IR Developer模式的發(fā)布，為開(kāi)發(fā)者打破了這堵“高墻”。

作為T(mén)ileLang-Ascend提供的兩大代碼生成后端之一（另一個(gè)是Ascend C），TileLang AscendNPU IR支持Expert和Developer開(kāi)發(fā)模式。如果說(shuō)此前推出的Expert模式是為追求100%極致性能的專家打造的“手術(shù)刀”，那么隨著Developer模式的演進(jìn)，TileLang-Ascend正基于AscendNPU IR強(qiáng)大的編譯推斷能力，致力于讓NPU算子開(kāi)發(fā)進(jìn)入“智能編譯”時(shí)代。

Developer模式：算子開(kāi)發(fā)的“智能編譯”

Developer模式的核心目標(biāo)，是實(shí)現(xiàn)從GPU到NPU跨架構(gòu)無(wú)縫遷移。對(duì)于習(xí)慣了GPU編程思維的開(kāi)發(fā)者，TileLang AscendNPU IR帶來(lái)了一個(gè)平滑的過(guò)渡路徑。

如下所示，在將一個(gè)經(jīng)典的矩陣乘（GEMM）算子從GPU遷移至?xí)N騰NPU時(shí)，開(kāi)發(fā)者面臨的代碼改動(dòng)微乎其微。得益于TileLang統(tǒng)一的T.Kernel和T.Parallel抽象，算子的核心切分邏輯（Tiling）和計(jì)算邏輯幾乎保持一致。開(kāi)發(fā)者僅需在編譯階段調(diào)整Backend Target，TileLang編譯器后端即可自動(dòng)將通用的并行原語(yǔ)映射為昇騰NPU特有的Cube/Vector指令，屏蔽了底層的架構(gòu)差異。

NPU算子代碼如下：

@tilelang.jit(out_idx=[-1], target="npuir")

def matmul(M, N, K, block_M, block_N, block_K, dtype="float16", accum_dtype="float32"):

@T.prim_func

def gemm(

A: T.Tensor((M, K), dtype),

B: T.Tensor((K, N), dtype),

C: T.Tensor((M, N), dtype),

with T.Kernel(T.ceildiv(N, block_N) * T.ceildiv(M, block_M), is_npu=True) as (cid, _):

by = cid // T.ceildiv(N, block_N)

bx = cid % T.ceildiv(N, block_N)

A_shared = T.alloc_shared((block_M, block_K), dtype)

B_shared = T.alloc_shared((block_K, block_N), dtype)

C_local = T.alloc_fragment((block_M, block_N), accum_dtype)

for k in T.Pipelined(T.ceildiv(K, block_K), num_stages=2):

T.copy(A[by * block_M, k * block_K], A_shared)

T.copy(B[k * block_K, bx * block_N], B_shared)

T.gemm(A_shared, B_shared, C_local, initC=(k == 0))

T.copy(C_local, C[by * block_M, bx * block_N])

return gemm

來(lái)源：https://github.com/tile-ai/tilelang-ascend/blob/npuir/examples/gemm/example_gemm.py

GPU算子代碼如下：

@tilelang.jit(out_idx=[-1])

def matmul(M, N, K, block_M, block_N, block_K, dtype=T.float16, accum_dtype=T.float32):

@T.prim_func

def gemm(

A: T.Tensor((M, K), dtype),

B: T.Tensor((K, N), dtype),

C: T.Tensor((M, N), dtype),

with T.Kernel(T.ceildiv(N, block_N), T.ceildiv(M, block_M), threads=128) as (bx, by):

A_shared = T.alloc_shared((block_M, block_K), dtype)

B_shared = T.alloc_shared((block_K, block_N), dtype)

C_local = T.alloc_fragment((block_M, block_N), accum_dtype)

T.clear(C_local)

for k in T.Pipelined(T.ceildiv(K, block_K), num_stages=3):

T.copy(A[by * block_M, k * block_K], A_shared)

T.copy(B[k * block_K, bx * block_N], B_shared)

T.gemm(A_shared, B_shared, C_local)

T.copy(C_local, C[by * block_M, bx * block_N])

return gemm

來(lái)源：https://github.com/tile-ai/tilelang/blob/main/examples/gemm/example_gemm.py

這充分體現(xiàn)了跨架構(gòu)無(wú)縫遷移的魅力——原有的Pythonic算子邏輯、Tiling策略以及循環(huán)結(jié)構(gòu)均可直接復(fù)用，開(kāi)發(fā)者無(wú)需重寫(xiě)核心算法，即可低成本完成跨硬件的算子遷移。

達(dá)成這一核心目標(biāo)的背后，是TileLang AscendNPU IR通過(guò)編譯器的智能分析與推斷，將原本需要開(kāi)發(fā)者手動(dòng)管理的底層硬件細(xì)節(jié)實(shí)現(xiàn)了“自動(dòng)化”。具體而言，Developer模式以如下四大核心特性，大幅降低了NPU編程的心智負(fù)擔(dān)：

1. 消除顯式CV分核，用戶零感知

在昇騰NPU的異構(gòu)架構(gòu)中，計(jì)算任務(wù)被分為Cube（矩陣計(jì)算）和Vector（向量計(jì)算）兩類核心。在傳統(tǒng)開(kāi)發(fā)模式下，開(kāi)發(fā)者必須顯式地指定哪段代碼運(yùn)行在Cube核，哪段運(yùn)行在 Vector核，并處理兩者間的交互。而在Developer模式下，編譯器具備了自動(dòng)推斷計(jì)算任務(wù)類型的能力。它能根據(jù)計(jì)算邏輯，自動(dòng)將任務(wù)分派給最合適的計(jì)算單元，開(kāi)發(fā)者無(wú)需再手動(dòng)編寫(xiě)T.Scope進(jìn)行分核指定，真正實(shí)現(xiàn)了對(duì)底層異構(gòu)核心的“零感知”。

2. 消除顯式內(nèi)存管理，自動(dòng)搬運(yùn)與同步

數(shù)據(jù)在Global Memory、L1、L0A/L0B、Unified Buffer (UB) 等多級(jí)存儲(chǔ)間流轉(zhuǎn)，是NPU編程中最繁瑣的部分。TileLang AscendNPU IR通過(guò)引入自動(dòng)內(nèi)存規(guī)劃，讓編譯器接管了這一切。這主要體現(xiàn)在如下兩方面：

跨硬件架構(gòu)的內(nèi)存抽象：通過(guò)T.alloc_shared、T.alloc_fragment等跨硬件架構(gòu)的兩層抽象，開(kāi)發(fā)者可以在不同的硬件內(nèi)存級(jí)別中靈活分配緩沖區(qū)。這既屏蔽了底層硬件的差異，又能確保性能的提升。

數(shù)據(jù)跨層自動(dòng)搬運(yùn)：利用T.copy接口，開(kāi)發(fā)者不再需要手動(dòng)管理復(fù)雜的跨層數(shù)據(jù)搬運(yùn)細(xì)節(jié)，編譯器會(huì)自動(dòng)處理數(shù)據(jù)在不同層級(jí)間的流動(dòng)，并插入必要的同步指令。

3. 自動(dòng)流水并行（T.Pipelined）

為了掩蓋內(nèi)存訪問(wèn)延遲，流水線并行（Pipeline Parallelism）是高性能算子的標(biāo)配。在Developer模式中，這一高階技巧被封裝得極為簡(jiǎn)潔。本質(zhì)上，通過(guò)T.Pipelined原語(yǔ)，編譯器能夠自動(dòng)實(shí)現(xiàn)Double Buffer和緩沖區(qū)重用。它不僅處理常規(guī)的數(shù)據(jù)依賴，更能自動(dòng)實(shí)現(xiàn)CV流水并行，將計(jì)算與數(shù)據(jù)搬運(yùn)重疊執(zhí)行，在不增加開(kāi)發(fā)者負(fù)擔(dān)的前提下，榨干硬件的每一分算力。

4. 向量化計(jì)算的自動(dòng)映射（T.Parallel）

針對(duì)Vector核心的計(jì)算，TileLang AscendNPU IR引入了T.Parallel原語(yǔ)，使得開(kāi)發(fā)者可以使用通用的數(shù)學(xué)符號(hào)（如a+b）或TileLang內(nèi)置函數(shù)（如T.exp, T.max）在T.Parallel作用域內(nèi)編寫(xiě)代碼，AscendNPU IR后端會(huì)自動(dòng)將其Lowering為昇騰硬件原生的向量指令，并自動(dòng)處理數(shù)據(jù)切分和Mask掩碼邏輯。這不僅保持了代碼的Pythonic風(fēng)格，更確保了在NPU Vector單元上的高效執(zhí)行。

實(shí)戰(zhàn)演示：CV融合算子的“智能編譯”

如下代碼展示了一個(gè)典型的CV融合算子開(kāi)發(fā)實(shí)例。在該場(chǎng)景中，我們需要先執(zhí)行矩陣乘法（Cube密集計(jì)算），隨后緊接一個(gè)Softmax（Vector向量計(jì)算）。

@tilelang.jit(target="npuir")

def minicv(M, N, K, block_M, block_N):

m_num = M // block_M

n_num = N // block_N

@T.prim_func

def main(

A: T.Tensor((M, K), dtype),

B: T.Tensor((K, N), dtype),

C: T.Tensor((M, N), inner_dtype),

D: T.Tensor((M, N), inner_dtype),

with T.Kernel(m_num * n_num, is_npu=True) as (cid, _):

blockx = cid // n_num

bx = blockx * block_M

blocky = cid % n_num

by = blocky * block_N

A_BUF = T.alloc_shared((block_M, K), dtype)

B_BUF = T.alloc_shared((K, block_N), dtype)

C_BUF = T.alloc_fragment((block_M, block_N), inner_dtype)

D_BUF = T.alloc_fragment((block_M, block_N), inner_dtype)

T.copy(A[bx, 0], A_BUF, [block_M, K])

T.copy(B[0, by], B_BUF, [K, block_N])

T.gemm(A_BUF, B_BUF, C_BUF, [block_M, K, block_N], initC = True)

T.exp (C_BUF, D_BUF)

T.copy(C_BUF, C[bx, by], [block_M, block_N])

T.copy(D_BUF, D[bx, by], [block_M, block_N])

return main

來(lái)源：https://github.com/tile-ai/tilelang-ascend/blob/npuir/examples/mixcv/example_mixcv.py

我們可以看到，在Developer模式下，代表Cube矩陣計(jì)算的T.gemm與代表Vector向量計(jì)算的T.exp被自然地書(shū)寫(xiě)在同一段邏輯中。

首先是零感知分核與同步：代碼中完全沒(méi)有顯式的SetFlag/WaitFlag同步原語(yǔ)，也無(wú)需手動(dòng)劃分Cube與Vector的執(zhí)行Scope，編譯器自動(dòng)識(shí)別出T.gemm歸屬Cube核，T.exp歸屬Vector核，并自動(dòng)插入同步指令處理數(shù)據(jù)依賴。

同時(shí)也實(shí)現(xiàn)了自動(dòng)流水與內(nèi)存管理：在整個(gè)計(jì)算過(guò)程中，開(kāi)發(fā)者無(wú)需管理復(fù)雜的L1/L0/UB內(nèi)存搬運(yùn)，編譯器后端自動(dòng)實(shí)現(xiàn)了Cube計(jì)算與Vector搬運(yùn)的流水線重疊（Double buffer）。

TileLang在昇騰生態(tài)中的獨(dú)特價(jià)值

在當(dāng)前的AI編譯器領(lǐng)域，TileLang-Ascend并非在重復(fù)造輪子，而是占據(jù)了一個(gè)不可替代的獨(dú)特生態(tài)位。與業(yè)界主流工具相比，它展現(xiàn)出了鮮明的差異化優(yōu)勢(shì)。

1. 硬件可感知，留住性能紅利

現(xiàn)有技術(shù)如Triton，其設(shè)計(jì)哲學(xué)是“硬件不可知”，試圖用一套邏輯屏蔽所有硬件細(xì)節(jié)。然而，這一注重通用性的技術(shù)路線，在擁有顯式存儲(chǔ)層級(jí)和存算分離（Cube/Vector）架構(gòu)的專用硬件上，往往難以觸達(dá)最底層的性能紅利。

TileLang AscendNPU IR則跑通了“硬件可感知”的路線。基于TileLang的硬件抽象能力，用戶可以通過(guò)注解輸入更多的信息（如T.alloc_shared等），這些信息幫助AscendNPU IR編譯器后端深度理解開(kāi)發(fā)者的意圖與NPU的硬件特性，從而做出更優(yōu)的指令調(diào)度與內(nèi)存分配。這種“懂硬件”的高級(jí)抽象，讓TileLang既能保持易用性，又能觸達(dá)通用編譯器難以觸及的底層優(yōu)化空間。

此外，由于AscendNPU IR是基于MLIR框架構(gòu)建的，它天然繼承了MLIR生態(tài)的靈活性。對(duì)于熟悉MLIR的開(kāi)發(fā)者而言，憑借既有經(jīng)驗(yàn)即可無(wú)縫上手，輕松擴(kuò)展自定義算子。同時(shí)，開(kāi)發(fā)者還可以根據(jù)需求，在AscendNPU IR編譯棧的任意層級(jí)進(jìn)行細(xì)粒度的優(yōu)化與調(diào)試。

2. 分層抽象，完善昇騰工具鏈

TileLang AscendNPU IR通過(guò)多級(jí)AscendNPU IR對(duì)接，構(gòu)建了靈活的開(kāi)發(fā)體系。它支持Developer和Expert等不同層級(jí)的開(kāi)發(fā)模式，既能滿足大眾開(kāi)發(fā)者對(duì)于算子編程易用性的訴求，又能滿足專家用戶對(duì)于極致性能的渴望，為昇騰生態(tài)提供了一個(gè)覆蓋全場(chǎng)景的算子開(kāi)發(fā)解決方案。

生態(tài)新范式：技術(shù)孵化與開(kāi)源共建

TileLang Ascend的快速成熟，不僅僅是一個(gè)技術(shù)項(xiàng)目的成功，更是昇騰“開(kāi)源開(kāi)放” 平臺(tái)戰(zhàn)略與“產(chǎn)學(xué)研協(xié)同”模式的縮影。

首先，該項(xiàng)目得到了“北京大學(xué)鯤鵬昇騰科教創(chuàng)新卓越中心”（簡(jiǎn)稱卓越中心）的強(qiáng)力支持。卓越中心致力于將高校在MLIR、編譯技術(shù)領(lǐng)域的前沿探索，迅速轉(zhuǎn)化為工業(yè)界的實(shí)戰(zhàn)工具，TileLang-Ascend正是這種“學(xué)術(shù)界前沿技術(shù)+產(chǎn)業(yè)界硬件底座”化學(xué)反應(yīng)的產(chǎn)物。

其次是開(kāi)源共建放大成果。昇騰作為一個(gè)開(kāi)放生態(tài)，AscendNPU IR為學(xué)術(shù)界及其他有創(chuàng)新需求的開(kāi)發(fā)者提供了基于昇騰平臺(tái)的創(chuàng)新研發(fā)機(jī)會(huì)，TileLang-Ascend從第一行代碼開(kāi)始就堅(jiān)持開(kāi)源，它不僅服務(wù)于華為內(nèi)部，更深度協(xié)同PyTorch、vLLM等開(kāi)源社區(qū)。

目前，社區(qū)已經(jīng)基于TileLang實(shí)現(xiàn)了FlashAttention、MatMul等核心算子，并正在吸引越來(lái)越多的開(kāi)發(fā)者參與共建，形成“技術(shù)孵化-開(kāi)源共建-生態(tài)反哺”的良性循環(huán)。

結(jié)語(yǔ)：讓算力觸手可及

隨著Developer模式基礎(chǔ)特性的陸續(xù)開(kāi)源，TileLang-Ascend正在將國(guó)產(chǎn)NPU算子開(kāi)發(fā)帶入一個(gè)自動(dòng)化、智能化的新階段。

對(duì)于開(kāi)發(fā)者而言，這意味著不再需要深陷于流水線編排和內(nèi)存管理的細(xì)節(jié)泥潭，只需關(guān)注算法邏輯本身。TileLang-Ascend正通過(guò)AscendNPU IR的技術(shù)突破，讓算力觸手可及。

想要了解技術(shù)路線的更多細(xì)節(jié)和代碼樣例的開(kāi)發(fā)者，請(qǐng)移步代碼倉(cāng)tile-ai/tilelang-ascend，關(guān)注npuir分支：https://github.com/tile-ai/tilelang-ascend/tree/npuir，共同見(jiàn)證這一技術(shù)變革。