近日，摩爾線程正式開(kāi)源TileLang-MUSA項(xiàng)目，實(shí)現(xiàn)對(duì)TileLang編程語(yǔ)言的完整支持。該項(xiàng)目已成功在摩爾線程多代全功能GPU上完成功能驗(yàn)證與特性開(kāi)發(fā)，旨在通過(guò)高層抽象與編譯器優(yōu)化，大幅降低開(kāi)發(fā)門檻，為國(guó)產(chǎn)算力平臺(tái)提供更高效的AI與高性能計(jì)算開(kāi)發(fā)體驗(yàn)。

▼TileLang-MUSA開(kāi)源地址：

https://github.com/MooreThreads/tilelang_musa

TileLang：重新定義GPU算子編程的語(yǔ)言

TileLang是一款基于張量分塊（Tiling）抽象的高性能AI算子編程語(yǔ)言，屬于領(lǐng)域特定語(yǔ)言（DSL）。它采用聲明式語(yǔ)法與類Python前端，使開(kāi)發(fā)者能夠以接近數(shù)學(xué)公式的形式描述計(jì)算意圖，并由編譯器自動(dòng)完成循環(huán)優(yōu)化、內(nèi)存調(diào)度與代碼生成，在保持底層性能的同時(shí)大幅降低GPU及異構(gòu)計(jì)算平臺(tái)的編程復(fù)雜度。

在實(shí)際應(yīng)用中，TileLang通過(guò)三大核心作用顯著提升GPU計(jì)算的開(kāi)發(fā)效率：

• 通過(guò)高級(jí)抽象降低開(kāi)發(fā)門檻，開(kāi)發(fā)者無(wú)需深入底層硬件知識(shí)即可生成高性能代碼；
• 具備跨平臺(tái)能力，實(shí)現(xiàn)“一次編寫、多架構(gòu)運(yùn)行”，有效解決多元算力生態(tài)的適配難題；
• 編譯器自動(dòng)執(zhí)行Layout推導(dǎo)、線程映射、Warp特化、流水線排布、內(nèi)存優(yōu)化等復(fù)雜優(yōu)化，在保障性能的同時(shí)提升開(kāi)發(fā)效率。

TileLang已廣泛應(yīng)用于多個(gè)關(guān)鍵領(lǐng)域，例如：在AI與機(jī)器學(xué)習(xí)中，研究人員可用其快速定義新型算子（如注意力機(jī)制）；在科學(xué)計(jì)算中，可便捷移植大型數(shù)值模擬程序至不同硬件平臺(tái)；對(duì)硬件廠商而言，TileLang可作為構(gòu)建芯片軟件生態(tài)的基礎(chǔ)工具鏈。

在產(chǎn)業(yè)實(shí)踐中，DeepSeek-V3的研發(fā)已采用TileLang進(jìn)行算子快速原型設(shè)計(jì)與性能驗(yàn)證，證明了其在大規(guī)模模型訓(xùn)練中的實(shí)戰(zhàn)價(jià)值。摩爾線程開(kāi)源的TileLang-MUSA項(xiàng)目，正是這一技術(shù)理念的產(chǎn)業(yè)落地——通過(guò)提供高效開(kāi)發(fā)工具鏈降低創(chuàng)新門檻，推動(dòng)國(guó)產(chǎn)算力應(yīng)用生態(tài)的繁榮發(fā)展。

TileLang-MUSA：連接前沿語(yǔ)法與國(guó)產(chǎn)算力的橋梁

摩爾線程此次開(kāi)源的 TileLang-MUSA項(xiàng)目，旨在充分釋放全功能GPU的性能潛力。它提供了一種介于底層匯編與高層DSL之間的“中間層”抽象，在保留硬件控制力的同時(shí)，顯著降低了編程復(fù)雜度。具體特性如下：

▼廣泛的硬件架構(gòu)覆蓋：TileLang-MUSA已在摩爾線程多代全功能GPU上完成功能驗(yàn)證與打通，包括訓(xùn)推一體全功能智算卡 MTT S5000和MTT S4000，展現(xiàn)了良好的硬件兼容性。

▼核心計(jì)算特性的深度映射：項(xiàng)目團(tuán)隊(duì)實(shí)現(xiàn)了TileLang高層語(yǔ)義到摩爾線程GPU底層MUSA架構(gòu)的精準(zhǔn)映射。

• Tensor Core 加速：編譯器能夠自動(dòng)調(diào)用MUSA的MMA（矩陣乘累加）指令，充分發(fā)揮硬件張量核心的峰值計(jì)算能力；
• Tile-Level Pipeline：自動(dòng)處理從全局內(nèi)存（Global Memory）到共享內(nèi)存（Shared Memory）再到寄存器（Registers）的多級(jí)數(shù)據(jù)搬運(yùn)，利用MUSA異步拷貝指令掩蓋訪存延遲；
• Warp級(jí)并行優(yōu)化：完整支持Warp Specialization特性。

目前，基于MUSA架構(gòu)的TileLang原生算子單元測(cè)試覆蓋率已超過(guò)80%，為大規(guī)模應(yīng)用提供了可靠保障。

代碼示例：體驗(yàn)“零門檻”算子遷移

開(kāi)發(fā)者在完成環(huán)境配置后，可保留原有的import tilelang習(xí)慣，通過(guò)Cython編譯后端直接在MUSA環(huán)境中運(yùn)行TileLang代碼。

以下是一個(gè)基于TileLang-MUSA的高性能矩陣乘法實(shí)現(xiàn)示例：通過(guò)簡(jiǎn)潔明了的語(yǔ)法描述計(jì)算邏輯，編譯器能夠自動(dòng)生成高度優(yōu)化的MUSA內(nèi)核代碼。依托TileLang豐富且經(jīng)過(guò)深度調(diào)優(yōu)的原語(yǔ)算子庫(kù)（primitives），開(kāi)發(fā)者借助TileLang-MUSA不僅能夠顯著提升編碼效率，更能直接調(diào)用底層硬件的計(jì)算潛力，實(shí)現(xiàn)媲美手工優(yōu)化性能的矩陣運(yùn)算。

import tilelang

import tilelang.language as T

import torch

from tilelang.primitives.gemm.base import GemmWarpPolicy

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

@T.prim_func

def matmul_kernel(

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=512) as (bx, by):

T.use_swizzle(panel_size=4, order='col')

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, policy=T.GemmWarpPolicy.Square)

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

return matmul_kernel

def main():

M, N, K, BLOCK_M, BLOCK_N, BLOCK_K = 8192, 7168, 16384, 256, 256, 64

device = "musa"

A = torch.randn((M, K), dtype=torch.float16, device=device)

B = torch.randn((K, N), dtype=torch.float16, device=device)

program = matmul(M, N, K, BLOCK_M, BLOCK_N, BLOCK_K, dtype="float16", accum_dtype="float32")

kernel = tilelang.compile(

program,

out_idx=-1,

target="musa",

execution_backend="cython",

verbose=True,

ref_out = torch.mm(A, B)

C = kernel(A, B)

torch.testing.assert_close(ref_out.to(torch.float16), C.to(torch.float16), rtol=1.25e-1, atol=1.25e-1)

if __name__ == "__main__":

main()

*代碼示例

實(shí)測(cè)表現(xiàn)：兼顧開(kāi)發(fā)效率與運(yùn)行性能

TileLang-MUSA的設(shè)計(jì)理念是讓開(kāi)發(fā)者“寫得快”且“跑得快”。在實(shí)際算子開(kāi)發(fā)實(shí)踐中，這一目標(biāo)已得到驗(yàn)證。

以大語(yǔ)言模型中至關(guān)重要的FlashAttention-3和GEMM（通用矩陣乘）算子為例，在摩爾線程MTT S5000上的測(cè)試結(jié)果顯示：

▼開(kāi)發(fā)效率倍增：相較手寫MUSA C++代碼，使用TileLang-MUSA的代碼量減少了約90%，且代碼邏輯更加清晰，極大降低了開(kāi)發(fā)與維護(hù)成本。

▼性能媲美手寫：得益于TileLang-MUSA優(yōu)秀的編譯器優(yōu)化，生成的算子性能在典型配置下，Gemm最高可達(dá)到手寫優(yōu)化版本的95%, FlashAttention-3可以達(dá)到手寫優(yōu)化版本的85%。

▼自動(dòng)化調(diào)優(yōu)：借助TileLang-MUSA的Auto-tuning機(jī)制，開(kāi)發(fā)者可以在MUSA架構(gòu)的全功能GPU上快速搜索最優(yōu)的分塊策略（Tile Size）和流水線級(jí)數(shù)，輕松超越未經(jīng)深度優(yōu)化的基準(zhǔn)實(shí)現(xiàn)。

TileLang-MUSA的推出，為開(kāi)發(fā)者帶來(lái)了三大核心價(jià)值：

▼無(wú)縫遷移：TileLang用戶可以近乎零成本地將算子邏輯遷移至摩爾線程GPU；

▼降低門檻：TileLang-MUSA為不熟悉MUSA指令集的AI工程師提供了高層次的開(kāi)發(fā)入口，大幅降低了國(guó)產(chǎn)GPU的應(yīng)用門檻；

▼賦能大模型：TileLang-MUSA通過(guò)支持FlashAttention等關(guān)鍵算子的高效開(kāi)發(fā)，將有力加速大語(yǔ)言模型等前沿AI應(yīng)用在國(guó)產(chǎn)算力平臺(tái)上的部署與落地。

TileLang-MUSA不僅驗(yàn)證了“高層語(yǔ)言+國(guó)產(chǎn)GPU”這一技術(shù)路線的可行性，更為構(gòu)建開(kāi)放、易用的國(guó)產(chǎn)算力開(kāi)發(fā)生態(tài)提供了關(guān)鍵工具支撐。

未來(lái)展望：打造基于MUSA的深度學(xué)習(xí)統(tǒng)一平臺(tái)

TileLang-MUSA的開(kāi)源是摩爾線程構(gòu)建國(guó)產(chǎn)算力生態(tài)的關(guān)鍵一步。以此為起點(diǎn)，摩爾線程將持續(xù)推進(jìn)平臺(tái)與生態(tài)建設(shè)，致力于打造一個(gè)覆蓋從單算子到完整大模型的國(guó)產(chǎn)算力統(tǒng)一加速平臺(tái)：持續(xù)進(jìn)行性能優(yōu)化，開(kāi)發(fā)更多MUSA架構(gòu)定制擴(kuò)展，使生成代碼性能穩(wěn)定達(dá)到手寫優(yōu)化版本的90%以上；深度集成SGLang等主流AI框架，實(shí)現(xiàn)訓(xùn)練與推理場(chǎng)景的端到端無(wú)縫加速；從單算子優(yōu)化延伸至Transformer、MoE等復(fù)雜模型架構(gòu)的跨算子調(diào)度與全局優(yōu)化；同時(shí)完善調(diào)試和性能分析工具鏈，為開(kāi)發(fā)者提供全流程支持。

這一切努力，最終都將匯聚于一個(gè)核心目標(biāo)：構(gòu)建一個(gè)開(kāi)放、高效、充滿生命力的國(guó)產(chǎn)算力生態(tài)，讓創(chuàng)新的想法在堅(jiān)實(shí)的基石上自由生長(zhǎng)。

了解TileLang-MUSA開(kāi)源項(xiàng)目可直接訪問(wèn)：

https://github.com/MooreThreads/tilelang_musa

*以上測(cè)試數(shù)據(jù)均來(lái)自摩爾線程實(shí)驗(yàn)室