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

TileLang-MUSA開源地址：

https://github.com/MooreThreads/tilelang_musa

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

TileLang是一款基于張量分塊（Tiling）抽象的高性能AI算子編程語(yǔ)言，屬于領(lǐng)域特定語(yǔ)言（DSL）。它采用聲明式語(yǔ)法與類Python前端，使開發(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ì)算的開發(fā)效率：

通過(guò)高級(jí)抽象降低開發(fā)門檻，開發(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í)提升開發(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à)值。摩爾線程開源的TileLang-MUSA項(xiàng)目，正是這一技術(shù)理念的產(chǎn)業(yè)落地——通過(guò)提供高效開發(fā)工具鏈降低創(chuàng)新門檻，推動(dòng)國(guó)產(chǎn)算力應(yīng)用生態(tài)的繁榮發(fā)展。

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

摩爾線程此次開源的 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)“零門檻”算子遷移

開發(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），開發(fā)者借助TileLang-MUSA不僅能夠顯著提升編碼效率，更能直接調(diào)用底層硬件的計(jì)算潛力，實(shí)現(xiàn)媲美手工優(yōu)化性能的矩陣運(yùn)算。

importtilelang
importtilelang.languageasT
importtorch
fromtilelang.primitives.gemm.baseimportGemmWarpPolicy


defmatmul(M, N, K, block_M, block_N, block_K, dtype="float16", accum_dtype="float"):
  @T.prim_func
defmatmul_kernel(
    A: T.Tensor((M, K), dtype),
    B: T.Tensor((K, N), dtype),
    C: T.Tensor((M, N), dtype),
):
withT.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)
forkinT.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])
returnmatmul_kernel


defmain():
  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)：兼顧開發(fā)效率與運(yùn)行性能

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

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

開發(fā)效率倍增：相較手寫MUSA C++代碼，使用TileLang-MUSA的代碼量減少了約90%，且代碼邏輯更加清晰，極大降低了開發(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ī)制，開發(fā)者可以在MUSA架構(gòu)的全功能GPU上快速搜索最優(yōu)的分塊策略（Tile Size）和流水線級(jí)數(shù)，輕松超越未經(jīng)深度優(yōu)化的基準(zhǔn)實(shí)現(xiàn)。

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

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

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

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

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

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

TileLang-MUSA的開源是摩爾線程構(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)化，開發(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)試和性能分析工具鏈，為開發(fā)者提供全流程支持。

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

點(diǎn)擊下方“閱讀原文”了解TileLang-MUSA開源項(xiàng)目，或直接訪問(wèn)：

https://github.com/MooreThreads/tilelang_musa

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

關(guān)于摩爾線程

摩爾線程以全功能GPU為核心，致力于向全球提供加速計(jì)算的基礎(chǔ)設(shè)施和一站式解決方案，為各行各業(yè)的數(shù)智化轉(zhuǎn)型提供強(qiáng)大的AI計(jì)算支持。

我們的目標(biāo)是成為具備國(guó)際競(jìng)爭(zhēng)力的GPU領(lǐng)軍企業(yè)，為融合人工智能和數(shù)字孿生的數(shù)智世界打造先進(jìn)的加速計(jì)算平臺(tái)。我們的愿景是為美好世界加速。