7 月 5 日，由 HyperAI 超神經主辦的 Meet AI Compiler 技術沙龍第 7 期如約而至。雖是盛夏酷暑，依舊擋不住大家的熱情——現場座無虛席，不少小伙伴甚至全程站著聽完每一場分享。來自 AMD 、沐曦集成電路、字節跳動、北京大學的多位講師輪番登臺，從底層編譯到實際落地，帶來了深刻的行業洞察與趨勢分析，干貨滿滿！

關注微信公眾號「HyperAI 超神經」，后臺回復關鍵字「0705 AI 編譯器」，即可獲取確認授權的講師演講 PPT 。

作為一種專為簡化高性能 GPU Kernel 開發而設計的編程語言，Triton 憑借簡化復雜并行計算編程的特性，已成為 LLM 推理與訓練框架中的關鍵工具。其核心優勢在于平衡開發效率與硬件性能——既避免了底層硬件細節的暴露，又能通過編譯器優化釋放 GPU 算力，這一特點使其在開源社區迅速普及。

身為 GPU 領域的領軍企業，AMD 率先實現了對 Triton 語言的支持，并將相關代碼貢獻至開源社區，推動 Triton 生態的跨廠商兼容。這一舉措不僅強化了 AMD 在高性能計算領域的技術影響力，更通過開源協作模式，為全球開發者提供了更靈活的 GPU 編程選擇，尤其是在大模型訓練與推理場景中，為算力優化提供了新路徑。

在以「助力開源社區，剖析 AMD Triton 編譯器」為題的演講中，來自 AMD 的 AI 架構師張寧圍繞公司在開源社區的技術貢獻，系統解讀了 AMD Triton 編譯器的核心技術、底層架構支撐及生態建設成果，為開發者深入理解高性能 GPU 編程與編譯器優化提供了全面視角。

HyperAI 超神經在不違原意的前提下，對張寧老師的演講分享進行了整理匯總，以下為演講實錄。

Triton：高效編程，實時編譯，靈活迭代

Triton 由 OpenAI 提出，是一種專為簡化高性能 GPU 內核開發設計的開源編程語言和編譯器，在主流 LLM 推理訓練框架中應用廣泛。其核心特性包括：

* 高效編程，能簡化 Kernel 開發，讓開發者無需深入了解復雜的 GPU 底層架構即可高效編寫 GPU 代碼；

* 實時編譯，支持即時編譯，可動態生成和優化 GPU 代碼以適應不同硬件與任務需求；

* 靈活的迭代空間結構，基于分塊程序和標量線程，增強了迭代空間靈活性，便于處理稀疏操作和優化數據局部性。

相比傳統方案，Triton 的優勢顯著：

首先，作為開源項目，Triton 提供基于 Python 的編程環境，用戶通過開發 Python Triton 代碼實現 GPU Kernel，無需關注底層 GPU 架構細節，大幅降低開發難度，相比 AMD HIP 等其他 GPU 編程方式能顯著提升產品開發效率。其編譯器會借助多種基于 GPU 架構特性的優化策略，將 Python 代碼轉換為優化后的 GPU 匯編代碼，實現上層張量操作到底層 GPU 指令的自動編譯，保障代碼在 GPU 上高效運行。

其次，Triton 具備良好的跨硬件兼容性，同一套代碼理論上可運行在英偉達、 AMD 的 GPU 以及支持 Triton 的國產 GPU 等多種硬件上。在性能與靈活性上，相比 PyTorch 等平臺能提供更好性能和優化靈活性，相比 CUDA 等則能夠隱藏底層 GPU 操作細節，讓開發者更專注于算法實現。

相較于 PyTorch API，更加專注于計算操作的具體實現，允許開發者靈活地定義線程塊切分方式，操作 Block/Tile 級別數據讀寫，并執行硬件相關的計算原語，特別適宜開發算子融合、參數調優等性能優化策略；

相較于 CUDA，Triton 隱藏了線程級別的操作控制，改由編譯器自動接管共享存儲、線程并行、合并訪存、張量布局等細節，降低了并行編程模型的難度，同時提高了 GPU 代碼的開發效率，達到開發效率和程序性能的有效均衡；開發者可以集中于算法的設計實現，而無需過度關心底層硬件細節和編程優化技巧，只要理解簡單的并行編程原理，即可快速開發性能較佳的 GPU 代碼。

從生態角度來看，它基于 Python 語言環境，使用 PyTorch 定義的張量數據類型，其函數可以無縫融入到 PyTorch 生態體系。相對于 CUDA 的封閉，Triton 的代碼開源和生態開放，也方便 AI 芯片廠商將其移植到自研芯片，利用開源社區來完善自己的工具鏈，同時也促進了 Triton 生態的健康發展。

AMD Triton 編譯器流程

Triton 編譯器包括 3 個關鍵模塊：前端模塊、優化器模塊和后端機器碼生成模塊，下圖所示：

Triton 編譯器模塊圖

前端模塊

前端模塊會遍歷 Python Triton 內核函數的抽象語法樹（AST），以創建 Triton 中間表示（Triton-IR），其內核函數會被轉換為 Triton-IR 。例如，標有 @triton.jit 裝飾器的 add_kernel 內核函數在這個模塊中會被轉成相應的 IR 。 JIT 裝飾器入口函數首先檢查 TRITON_INTERPRET 環境變量的值，如果該變量為 True，則調用 InterpretedFunction，以解釋模式運行 Triton 內核；如果為 False，則運行 JITFunction，在實際設備上編譯并執行 Triton 內核。

內核編譯的入口點是 Triton compile 函數，該函數會與目標設備及編譯選項信息一起調用。此過程將創建內核緩存管理器，啟動編譯流水線，并填充內核元數據。此外，它還加載特定于后端的方言，例如用于 AMD 平臺的 TritonAMDGPUDialect，以及處理 LLVM-IR 編譯的后端專用 LLVM 模塊。如果所有準備工作就緒，將調用 ast_to_ttir 函數生成該內核的 Triton-IR 文件。

優化器模塊

優化器模塊分為 3 個核心部分：Triton-IR 優化、 Triton-GPU IR 優化和 LLVM-IR 優化。

* Triton-IR 優化

在 AMD 平臺上，Triton-IR 優化流程由 make_ttir 函數定義。在這一階段，優化過程與硬件無關，包括內聯優化、通用子表達式消除、規范化、死代碼消除、循環不變代碼移動以及循環展開。

* Triton-GPU IR 優化

在 AMD 平臺上，Triton-IR 優化流程由 make_ttgir 函數定義，用于提升 GPU 性能。基于 AMD GPU 特點和優化經驗，我們還開發了特定的優化流程，如下圖所示:

AMD GPU-IR 特定優化

首先，針對 AMD GPU 的加速矩陣乘法，優化重點在于 Triton 中點乘操作的輸入/輸出布局，使其更好地兼容 AMD 的矩陣計算單元（如 Matrix）。該優化涉及大量針對 CDNA 架構的匹配處理，是整個實現中工作量較大的一部分。如果希望在 AMD 平臺上對 Triton 生成的代碼進行深度優化，這部分實現值得重點參考。

其次，在尾部處理階段，優化策略為直接存儲累加器，從而避免使用共享內存，減少共享內存的訪問壓力，提高整體效率。

接著，在 GPU 流計算過程中，引入了加載與計算的流水線化處理。即在前一個任務執行的同時，調用相應的內存進行數據加載，形成加載與計算并行執行的模式。這種機制在許多用戶場景中已有良好效果。在流水線優化基礎上，還引入了指令調度提示，用于在計算單元或遠程操作完成后對指令流進行引導，提升指令級別的響應效率。

隨后，AMD 實現了多套針對不同目標的指令重排序，包括：緩解寄存器壓力、避免冗余資源分配與釋放、優化加載-計算流程的銜接等。其中一部分重排序與流水線機制緊密結合，另一部分則聚焦于調整 LLVM IR 的指令順序，以更好地服務于 AMDGCN 匯編的生成規則。

除了指令重排序與調度提示外，我們還引入了另一項調度優化策略——乒乓調度。通過循環調度機制，讓兩個線程束在同一個 SIMD 單元上交替執行，避免空閑與等待，從而提高計算資源的利用率。

此外，AMD 還進行了指針規范化與緩沖區操作轉換方面的優化。該優化的主要目標是將指令有效映射到具體業務應用中，實現更高效的原子指令執行。

這些優化流程首先將 Triton-IR 轉換為 Triton-GPU IR 。在此過程中，布局信息被嵌入到 IR 中。以下圖為例，張量以 #blocked 布局形式表示。

流程示例圖

如果嘗試另一個 Triton 矩陣乘法示例，通過上述優化流程引入共享內存訪問以提高性能，這是矩陣乘法的常見優化解決方案，同時使用了設計用于 AMD MFMA  加速器的 amd_mfma 布局。

最后，在實驗驗證中，我以一個復雜度更高的 Matrix Multiplication 矩陣層為例。在確認 GPR（General Purpose Register）映射的過程中，插入了大量硬件相關指令，如 MFMA 、轉置（Transposed）、共享內存調用等。通過減少 bank 沖突、應用 swizzled 操作等方式，結合指令重排序與流水線化處理策略，所有這些優化都會在 Triton IR 轉換過程中自動加入，實現更強的硬件適配能力與性能提升。

流程示例圖

* LLVM-IR 優化

在 AMD 平臺上，優化流程通過 make_llir 函數定義。此函數包括 2 個部分：IR 級別優化和 AMD GPU LLVM 編譯器配置。對于 IR 級別優化，AMD GPU 特定的優化流程包括 LDS/共享內存相關優化和 LLVM-IR 級別的通用優化，如下圖所示:

LLVM-IR 特定優化流程

首先，對部分 AMD 不支持的轉換操作進行分解處理。例如，在 Triton GPU IR 向 LLVM-IR 的轉換過程中，如果遇到當前尚不支持的轉換路徑，我們會將這些操作拆解為更基礎的子操作，從而確保整個轉換流程能夠順利完成。

在配置 AMD GPU LLVM 編譯器時，首先初始化 LLVM 目標庫和上下文，設置 LLVM 模塊上的編譯參數，然后為 AMD GPU HIP 內核設定調用約定，并配置一些 LLVM-IR 屬性，如 amdgpu-flat-work-group-size 、 amdgpu-waves-per-eu 和 denormal-fp-math-f32，最后運行 LLVM 優化，并將優化級別設為 OPTIMIZE_O3 。

屬性配置參考文檔：

??https://llvm.org/docs/AMDGPUUsage.html??

后端機器代碼生成模塊

后端機器代碼生成模塊主要負責將中間代碼轉化為可在硬件上運行的二進制文件。該階段主要分為 2 個步驟：生成 AMDGCN 匯編代碼和構建最終的 AMD hsaco ELF 文件。

首先調用 translateLLVMIRToASM 函數，在 make_amdgcn 階段生成 AMD 匯編代碼。這一過程完成了從中間代碼到目標架構指令集的映射，為后續的二進制生成奠定了基礎。隨后，編譯器通過 assemble_amdgcn 函數，配合 ROCm 的鏈接模塊，在 make_hsaco 階段生成 AMD hsaco ELF（Executable and Linkable Format）二進制文件，該文件即為可在 AMD GPU 上直接運行的最終二進制，包含了完整的設備端指令與元信息。

通過這兩個步驟，編譯器將高級中間表示高效轉換為底層 GPU 可執行代碼，確保了程序能夠在 AMD Instinct 系列等 GPU 上順利運行，并充分發揮底層硬件性能。

AMD GPU 開發者云

AMD 正式面向全球開發者和開源社區開放其高性能 GPU 云平臺 —— AMD Developer Cloud，旨在讓每一位開發者都能無障礙接入世界級計算資源，便捷地訪問 AMD Instinct MI 系列 GPU 資源，快速上手 AI 和高性能計算任務。

在 AMD Developer Cloud 中，開發者可以根據需求靈活選擇計算資源：

* 小型：1 個 MI 系列 GPU（192 GB 顯存）

* 大型：8 個 MI 系列 GPU（1536 GB 顯存）

平臺最大程度降低了配置門檻，用戶無需繁瑣安裝，可即時啟動基于云的 Jupyter Notebook 。只需一個 GitHub 賬戶或郵箱，即可輕松完成配置。此外，AMD Developer Cloud 提供了預配置的 Docker 容器，內置主流 AI 軟件框架，最大限度減少了環境搭建時間，同時也保留了高度靈活性，允許開發者根據具體項目需求自定義代碼。

歡迎各位開發者來親自體驗 AMD Developer Cloud，在這里運行你的代碼、驗證你的想法，平臺將為你提供穩定、強大、靈活的算力支持，加速創新與落地。

AMD Developer Cloud 鏈接：

??https://www.amd.com/en/developer/resources/cloud-access/amd-developer-cloud.html??

獲取 PPT：關注微信公眾號「HyperAI 超神經」，后臺回復關鍵字「0705 AI 編譯器」，即可獲取確認授權的講師演講 PPT 。