網易首頁 > 網易號 > 正文申請入駐

AMD GPU性能暴漲7倍，優化算法首次開源！高效MoE支持任意專家數量

2025-04-03 13:27:05　來源: 新智元

北京舉報

分享至

新智元報道

編輯：LRST

【新智元導讀】通過完全啟用并發多塊執行，支持任意專家數量（MAX_EXPERT_NUMBER==256），并積極利用共享內存（5kB LDS）和寄存器（52 VGPRs，48 SGPRs），MoE Align & Sort邏輯被精心設計，實現了顯著的性能提升：A100提升3倍，H200提升3倍，MI100提升10倍，MI300X/MI300A提升7倍...

MoE（Mixture of Experts）模型模仿了人腦的低功耗運作模式：功能被劃分為多個獨立的部分，在思考時通過自適應路由部分激活，從而提高計算效率。

牛津大學研究論文中的人腦皮層示意圖，來源于互聯網

首個可在CUDA真正可行的版本是Switch Transformer[1]，隨后通過循環利用(Up Cycling)稠密模型Mistral[2]進一步優化了該設計。

SwitchTransformer-MoE

隨后，DeepSeek V2/V3/R1[3][4][5]通過引入共享專家[3]和門控偏差（gating bias）[4][5]進一步改進了MoE，最終實現了無輔助損失（auxiliary loss free）的MoE模型 [4][5]。這一優化本質上歸因于一個關鍵事實：當使用共享專家（DeepSeek團隊選擇的值為1）時，可以通過在較大的專家池（256個上施加偏差分數的懲罰，從而緩解專家路由的不均衡問題[11]。

MoE層本質上是由多個專家前饋網絡（FFN）組成的層，其中包含門控函數（gating functions），用于根據Top-K門控分數（DeepSeek V3/R1中引入偏差）進行激活路由，并在所選的FFN層上通過Group GEMM計算logits。

該功能在很大程度上依賴于基數排序（radix sort）邏輯。借助MoE Align & Sort，機器學習研究人員和實踐者可以按照專家ID對tokens進行排序。

在某些應用中，例如TransformerEngine[6][7]，該操作最初是通過已廢棄的cub::DeviceRadixSort實現的，而增加的permute操作用于記錄源（左）到目標（右）的映射，其梯度操作為unpermute。

MoE Permute示例

盡管cub::DeviceRadixSort大量使用共享內存，相比于基于__shfl_xor_sync（僅使用線程本地內存）的實現略慢，但它不支持對齊排序（alignment sorting）。

對齊排序對于Group GEMM的效率至關重要，因為它允許專家以塊（block 為單位處理tokens。

SGLang 中的MoE Align & Sort算法采用了對齊排序，但在支持多達256個專家的大規模prefill操作時效率并不理想。該問題已在issue#2732中被確認。

目前的實現將MoE Align & Sort拆分為兩個kernel啟動（kernel launches）：

對齊（alignment）：在單個block內執行傳統基數排序算法對齊后的偏移計算（alignment-based offsets computation）;
放置（placement）：根據在多個block并行計算出的偏移量，并行放置tokens;

研究人員提出并編寫了AMD友好的CUDA設備代碼，采用了該設計的MoE Align & Sort算法。因此，在AMD平臺上的性能分析和優化將被充分考慮。

文章地址：https://shorturl.at/C23JF

通過在不同的工作負載下使用RocProfiler-Compute進行分析，可以清楚地看到，即使不計入多次設備函數啟動的額外開銷，第一個kernel仍然消耗了33W個周期，第二個kernel消耗了8W個周期，總計41W周期：

the moe align kernel 1

the moe align kernel 2

在ROCm SDK 6.3.0 中，omniperf已更名為rocprof-compute。盡管MI300X/MI300A已得到積極支持，但該工具默認未隨ROCm SDK 6.3.0一同發布。不過，在Tools-dockerhub中的展示一樣，ROCm計算分析工具的設置僅需簡單三步。

現在，在PR#3613（https://shorturl.at/OPbiI）中應用優化方案后，片上計算開銷將從之前的41W個周期立即降低至20W個周期。

在SGLang中實現高效的多塊（multi-blocks）MoE-Align

通過完全地多塊（multiple blocks）并發執行，并支持任意專家數量（MAX_EXPERT_NUMBER==256），結合激進使用共享內存（5kB LDS）和寄存器（52 VGPRs，48 SGPRs），MoE Align & Sort邏輯被優化，實現了以下性能提升3x in A100,3x in H200, 10x in MI100, and 7x in MI300X/Mi300A:

借助Rocprof-Compute，可以輕松收集捕獲設備代碼的一些關鍵指標，并在遠程GUI服務器上進行可視化展示：

服務端開啟Rocprof-Compute

總而言之，在AMD MI300X/MI300A上，所提出的高效多塊（multi-blocks）MoE Align & Sort算法充分利用了每個wave的向量寄存器（52個），且無寄存器溢出（我已將初始線程塊大小調整至最佳值）；同時，每個CU使用5kB LDS，且僅有6.8%的存儲銀行沖突率。

研究人員還分析了MoE Sort & Align的Roofline模型。該模型顯示，設備代碼的性能在受限于內存帶寬的區域有所下降。

在AMD Compute Profile部分，研究人員詳細介紹了在ROCm平臺上算法設計的影響與性能數據。

本質上，MI300X/MI300A是全球首款基于多芯片（multi-die）設計的高性能AI加速器架構。因此，在該芯片上進行算子優化的方式將與NVIDIA平臺略有不同。

基本規則是，XCDs（加速計算芯片）之間的同步代價較高，因此最好充分利用XCDs，并利用L2緩存的局部性親和性來提高性能。

此外，研究人員應避免昂貴的同步開銷，具體方法包括：

當網格大小小于每顆芯片上的XCD數量（MI300X為8，MI300A為6）時，優先使用最低速計算單元（MI300X使用XCD7，MI300A使用XCD5）。
當網格大小大于每顆芯片上的XCD數量時，將其調整為XCD數量的整數倍。

使用hipCooperativeLaunch啟動協作設備代碼可能會增加L2緩存壓力（與紋理尋址器停滯率和忙碌率相關），特別是在數據交換（尤其是Die-Die交換增多的情況下。

在此示例中，之前main分支的實現使用了39個活躍CU，這已經接近最佳，因為本質上使用了兩個Die。

該實現在多塊（multi-blocks）執行中使用了66個活躍CU，跨越兩個Die，并且塊級歸約（block-wise reduction）過程中Die-Die數據交換是不可避免的，將在本季度晚些時候向SGLang提交進一步的V4優化。

具體細節將在性能分析（profiling）部分進一步討論。

SGLang中Fused MoE的回顧

SGLang團隊采用Triton First方法實現了相關邏輯，并在2024年12月成功實現DeepSeek V3的Day-0支持。

SGLang的MoE調用了使用Triton實現的Fused MoE 設備代碼。

在設備代碼啟動之前，會應用MoE Align & Sort算法。MoE Align & Sort的Triton設備代碼被拆分為四個階段，其中直接訪問DRAM，而不使用共享內存，這與向量化 Triton版本形成對比。

與單塊（single block wise）CUDA實現相比，Triton版本的多次設備代碼觸發以及對LDS、本地緩存和寄存器（例如VGPR）的低效利用，導致了在小規模工作負載上的單次測試執行效率較低。

隨后，CUDA實現最終被拆分為兩個階段，其中僅第二階段的執行在多塊（multiple blocks）上進行了加速。

MoE Align & Sort CUDA算法在其他開源平臺的實現

FasterTransfomer

在Mistral[2]和DeepSeek V2[3]之前，開放式稠密模型（open dense models）在推理場景中更為流行。這也是FasterTransformer[8]誕生的時期。

在FasterTransformer[8]項目中（由NVIDIA發起），MoE模型的支持主要依賴于cub::DeviceRadixSort，以及諸如moe_softmax（本質上是cub::BlockReduce實現的softmax）、moe_top_k及其融合版本topk_gating_softmax、用于排列潛在向量logits的permute，最終執行group gemm。

因此，融合優化主要（按計算開銷計算）限制在topk gating softmax和biased topk gating softmax，后續這些優化被整合進SGLang。

Megatron

在該算法發表之前，Megatron在FP16/BF16計算中主要采用FasterTransformer方法，但額外添加了permute的梯度操作unpermute，以支持訓練任務。

這意味著MoE仍然沒有得到高效融合。

vLLM

SGLang使用了許多vLLM設備代碼，但vLLM的Fused MoE最初是由SGLang團隊貢獻的。因此，它們采用了相同的方法進行部署。

首個AMD友好的Fused MoE版本于2024年11月26日在CK#1634（https://tinyurl.com/3fuj7yws）中提出。隨后，MoE Align & Sort被添加到CK#1771（https://tinyurl.com/5h4e8jat）和CK#1840（https://tinyurl.com/3wm8pdc3）中。

核心思路是將MoE 排序與Group GEMM進行融合。此外，CK中的MoE & Sorting在很大程度上采用了SGLang團隊的方法，但在CK pipeline及partitioner方面有所不同。

CK融合MoE思路[9]

融合per_group_token_quant（用于在線FP8量化）、MoE排序和Group GEMM可以通過將Radix Sort計算邏輯納入Group GEMM pipeline輕松解決：即統計出現次數以計算偏移量，隨后進行并行放置。

其中最關鍵的問題之一是如何平衡Radix Sorting和Group GEMM這兩種計算負載。

在AMD數據中心芯片中，Group GEMM片段更可能均勻分布在XCD內的所有可用計算單元。然而，當涉及多個XCD時，不同CU之間的數據交換主要通過低速L2 Cache及其互聯結構（L2 Cache fabric）進行。

編寫CK設備代碼需要先編寫主機端CK解決方案啟動器：

// Here is the entry of fused MoE :
//   https://github.com/ROCm/composable_kernel/blob/1342ecf7fbf64f43d8621cf6665c583fdc49b2c6/example/ck_tile/15_fused_moe/instances/fused_moegemm_api_internal.hpp
using f_pipeline    = ck_tile::FusedMoeGemmPipeline_FlatmmUk ;
using f_partitioner = ck_tile::FusedMoeGemmTilePartitioner_Linear ;
using f_kernel      = ck_tile::FusedMoeGemmKernel void >;
const dim3 grids                       = f_kernel::GridSize(a);
constexpr dim3 blocks                  = f_kernel::BlockSize();
constexpr ck_tile::index_t kBlockPerCu = 1;
static int printed = 0;
auto kargs = f_kernel::MakeKargs(a);
if(s.log_level_ > 0 && printed == 0)
{
std::cout << ", " << f_kernel::GetName() << std::flush;
printed = 1;
}
return ck_tile::launch_kernel(
s, ck_tile::make_kernel (f_kernel{}, grids, blocks,  0, kargs));

AMD CK分區器和階段流水線（stages pipeliner）在Fused MoE的最終匯編過程中扮演了重要角色，確實值得深入研究，但已超出本文討論范圍。

但需要記住，MoE Align & Sort是生產者代碼的一部分：

// https://github.com/ROCm/composable_kernel/blame/fdaff5603ebae7f8eddd070fcc02941d84f20538/include/ck_tile/ops/fused_moe/kernel/moe_sorting_kernel.hpp#L438
CK_TILE_DEVICE void moe_align_block_size_kernel(...)
{
const index_t tid       = static_cast
             
 (threadIdx.x); const index_t start_idx = tid * tokens_per_thread; ... #if 1 if(tid < num_experts){ // each thread reduce a column segment of tokens_cnts with # blockDim.x elements ... } #else ... #endif __syncthreads(); // do cumsum to compute offsets based on condition ... // do parallel placement based on the offsets computed ... }

因此，在AMD CK方案中，MoE Align & Sort的實現幾乎與SGLang主實現保持一致，僅在分區器（partitioner）和流水線（pipeliner）方面有所不同。

需要注意的是，該實現并不總是能在AMD平臺上提供最佳性能（請參考AITER中的asm MoE）。

由于AMD CDNA3架構并不支持類似Graphcore的片上（on-chip）洗牌操作（在2023年已經將PopART[12] & PopRT的Remapping操作進行抽象與泛化），而這一特性已在NVIDIA H100/H200/B200中得到了支持，并通過高效的SM<->SM片上通信實現。

因此，在AMD 開源解決方案中，如何以低開銷方式在塊（block）之間優化數據布局將是一個非常有趣的研究方向。

從哲學上講，這兩類不同工作負載的基于 Tiling 的融合代碼可能并不總是比非融合版本更優。相關研究的詳細內容將在V4 版本發布時進一步探討。

AITER

AI Tensor Engine[10]

AITER在今年早些時候被引入，以便整合在不同項目中使用的LLM設備代碼。它通過ck moe、asm 版本的 MoE 通過 hipModule和triton fused moe支持MoE融合。

因此，AITER是部分開源的，因為不透明的匯編代碼和開發計劃是針對MI300X開發者的。

AITER中fused MoE的三倍加速[10]已由Bruce Xu[13]驗證，并且這一加速主要來自于在不同形狀的Group GEMM中觀察到的加速：一個GEMM操作，其中每個專家的FFN權重與一塊隱藏狀態的token進行相乘。

這一證明可以在PR#199（https://shorturl.at/F8y0F）中找到，asm gemm幾乎帶來了三倍的性能提升。

ASM版本扁平矩陣乘

值得注意的是，仍然有一些情況下，選擇了來自SGLang社區的triton設備代碼。為了在MI300X/MI300A上高效運行triton設備代碼，它們采用了基于多芯片架構的特定邏輯，將線程塊映射到不同的計算單元（dies）上：

# https://github.com/ROCm/triton/blob/f669d3038f4c03ee7a60835e875937c65b5cec35/python/perf-kernels/gemm.py#L115
...
## pid remapping on xcds
# Number of pids per XCD in the new arrangement
pids_per_xcd = (GRID_MN + NUM_XCDS - 1) // NUM_XCDS
# When GRID_MN cannot divide NUM_XCDS, some xcds will have
# pids_per_xcd pids, the other will have pids_per_xcd - 1 pids.
# We calculate the number of xcds that have pids_per_xcd pids as
# tall_xcds
tall_xcds = GRID_MN % NUM_XCDS
tall_xcds = NUM_XCDS if tall_xcds == 0 else tall_xcds
# Compute current XCD and local pid within the XCD
xcd = pid % NUM_XCDS
local_pid = pid // NUM_XCDS
# Calculate new pid based on the new grouping
# Note that we need to consider the following two cases:
# 1. the current pid is on a tall xcd
# 2. the current pid is on a short xcd
if xcd < tall_xcds:
pid = xcd * pids_per_xcd + local_pid
else:
pid = tall_xcds * pids_per_xcd + (xcd - tall_xcds) * (pids_per_xcd - 1) + local_pid
if GROUP_SIZE_M == 1:
pid_m = pid // num_pid_n
pid_n = pid % num_pid_n
else:
num_pid_in_group = GROUP_SIZE_M * num_pid_n
group_id = pid // num_pid_in_group
first_pid_m = group_id * GROUP_SIZE_M
group_size_m = min(num_pid_m - first_pid_m, GROUP_SIZE_M)
pid_m = first_pid_m + (pid % group_size_m)
pid_n = (pid % num_pid_in_group) // group_size_m
...

此外，在CK fused MoE中使用了多種AMD芯片內建函數（intrinsics），例如：

__builtin_nontemporal_load,
__builtin_amdgcn_ds_swizzle,
__builtin_amdgcn_ds_permute/__builtin_amdgcn_ds_bpermute,
_builtin_amdgcn_mov_dpp

等等。這些內建函數可能最終影響fused MoE的匯編實現和性能。

例如，使用__builtin_nontemporal_load可以跳過L2緩存，從而為預測將被重復使用的數據留出更多L2緩存行空間。

Cutlass v3.8

Fused MoE尚未在NVIDIA Cutlass 3.8.0中公開支持。因此，當前該倉庫中沒有提供MoE Align & Sort功能。

TRT-LLM

在v0.16.0之前，TRT-LLM基本上遵循了FasterTransformer的方法。自v0.17.0版本起，MoE部分開始公開。

編寫對AMD設備友好的CUDA實現，并帶來超過3x~7x加速

該算法采用了多塊執行方案，并由三個不同的部分（D-C-P）組成：

分布式并發計數
計算累積和（cumsum）
- 并行非對齊本地累積和
- 合并非對齊累積和
- 對齊全局累積和
- 存儲全局累積和
并行放置

高效MoE Align& Sort算法

并行非對齊本地累積和

該算法首次由在PR#2970（https://shorturl.at/CuBs5）中提出并實現。

研究人員將每個塊中的累積和執行進行了負載均衡，分配給kElementsPerThr(16)個線程，每個線程需要處理kElementsPerThr+kElementsPerThr+threadIdx.x次加法操作。

因此，與當前倉庫中的單線程版本相比，波前（wavefront）更快地到達，可以觀察到此版本實現的性能提升了30%。

合并非對齊累積和

一旦獲得了每個塊中的本地非對齊累積和（Unaligned Cumsum），就可以在預分配的HBM緩沖區中進行塊級別的累積和歸約。

研究人員選擇了FRAG_SIZE_M(16)xFRAG_SIZE_N(16)xFRAGS_PER_BLOCK(4)的SRAM塊進行塊級歸約，其中FRAGS_PER_BLOCK是可調的：

塊級歸約

在AMD平臺上，計算是基于「1 warp加載/1warp計算」的方式進行的，而在NVIDIA平臺上則是「2warps加載和1warp計算」。

該設計充分利用了AMD CDNA3架構中64個SIMD通道的優勢。并且，在這種多芯片架構中，塊的數量始終是XCD數量的倍數。

FRAGS_PER_BLOCK被設置為4，以便在多輪中復用SMEM。

對齊全局累積和和存儲全局累積和

研究人員改進了向量化代碼，并處理了如果輸入數據大小與kElementsPerAccess常量不對齊時的循環尾部情況。

基準測試顯示，合并率有所提高，但仍然限制在30%左右。研究人員將在V4版本中繼續優化此問題。

編寫AMD友好的CUDA代碼

編寫PyTorch擴展可以自動將CUDA設備代碼轉換為HIP設備代碼，配合ROCm SDK進行使用。

但是，有些情況下HIP設備代碼與CUDA設備代碼表現不同：

Warp大小是一個與架構相關的全局變量，并在ROCm SDK中定義為warpSize；在CDNA3架構中，warpSize定義為64。
設備函數簽名可能與CUDA不完全對齊，因此需要條件編譯來支持這些符號。
需要特別關注多芯片架構中的L2緩存優化。

基準測試

在沒有CUDA圖捕獲的情況下，研究人員針對DeepSeek V3模型的大規模工作負載進行了廣泛測試。因此，專家數量設置為256。當前的算法不支持在CUDA圖捕獲下運行，將在V4版本中解決此問題。

由于GPU虛擬化和測試節點上分配的CPU數量，性能可能會與裸機測試時有所不同。

因此，研究人員使用Triton實現作為基準，展示MoE Align & Sort算法在加速倍數和效率上的表現。

每個測試首先進行了驗證，之后才開始基準測試。在基準測試中，可以觀察到，在AMD平臺上，Triton的運行時間顯著長于在NVIDIA平臺上的運行時間，因此建議進一步優化Triton的MLIR，以獲得比NVIDIA Triton更高效的降級過程。

對于AMD Triton，可以觀察到MI300X的速度比MI100快1.5倍，因此MI300X的性能提升幅度不像MI100那么顯著。此外，盡管普遍認為MI300X比MI100更快，但在測試中，MI100上的算法性能要優于MI300X。

這部分歸因于內存瓶頸操作，在多芯片之間的通信降低了執行速度。

在兩個平臺上，都觀察到了應用該算法后顯著的性能改進，其中現有的CUDA實現幾乎與Triton消耗相同的時間。

AMD系統準備

為了最大化使用AMD異構系統，建議進行以下檢查。

NVIDIA Grace CPU和AMD EPYC 9004系統通常建議禁用NUMA自動平衡，以便與GPU協同工作；然而，在某些情況下，可能不建議禁用 NUMA自動平衡。
啟用虛擬化時，建議啟用IOMMU直通模式，以消除DMA翻譯，從而帶來性能提升。

MI100基準測試

git clone 
https://github.com/yiakwy-xpu-ml-framework-team/AMD-sglang-benchmark-fork.git
 -b optimize_moe_align_v3 && cd sgl-kernel && python setup_rocm.py install

可以驗證不同輸入令牌和專家數量組合的可行性 :

cd ../benchmark/kernels/fused_moe_trition && python benchmark_deepseekv3_moe_align_blocks.py --verify

A100 性能測試

H200 性能測試

MI300X 性能測試

AMD Compute Profile

設置

在ROCm 6.3.3版本中，設置rocprof-compute只需三步即可完成，詳細的設置步驟可以在這里找到：Tools-dockerhub中的rocprof-compute設置。

向量L1緩存的分析結果

在分析中，工作負載為16384個tokens x（從256個專家中選擇8個），除非另有說明。

研究人員在算法中最大化了VGPRs的使用，但減少了SGPRs的總使用量。數據也表明，VGPRs/SGPRs的溢出為零，這表明寄存器的使用是健康的，并且此設備代碼沒有性能損失。

向量L1緩存（vL1D）是每個CU的本地單元，命中率記錄了從L2緩存請求到CU時的緩存行命中率。30%的L2緩存請求通過vL1D的紋理尋址器合并，達到了61%的命中率，如果需要，稍后可以進一步提升。

當數據從CU請求到vL1D的尋址處理單元（紋理尋址器）時，復雜的決策邏輯決定是否接受數據請求或回滾數據請求。以下是四種狀態：

Busy（忙碌）：紋理尋址器正在處理地址。
Address Stall（地址停頓）：紋理尋址器無法發送地址到vL1D。
Data Sending Stall（數據發送停頓）：紋理尋址器無法發送數據到vL1D。
Data Waiting Stall（數據等待停頓）：紋理尋址器等待發送數據到vL1D的數據處理單元。

有關這種微架構行為的詳細信息，可以在AMD CDNA3的ISA文檔以及rocProfiler-compute文檔中找到。

vL1D 尋址器停頓

研究人員在該算法設計中觀察到了18.61%的數據等待停頓率來自于向量L1緩存。

數據的讀寫負載平衡大大減少，從8kB的讀取操作和27B的寫入操作，轉變為109B的讀取操作，468B的寫入操作和202B的原子操作的組合。

L2緩存的分析結果

在CDNA3架構中，L2緩存是所有計算單元（CU）共享的，且是線程塊之間共享數據的主要通道，這些線程塊分布在不同的CUs上。

通過多通道和地址交錯設計，向L2緩存的請求可以大大并行處理。

此外，使用AMD特有的內置函數如__builtin_nontemporal_load，可以繞過L2緩存來處理那些不需要再次訪問的數據。

更多L2緩存研究細節將在V4版本中揭示。

結論

新的算法通過最大化使用LDS和向量寄存器，顯著加速了CUDA和ROCm平臺上的MoE Align & Sort，提升幅度高達3x~7x。

還可以觀察到，相較于單個芯片，內存密集型操作在多晶粒異構集成架構下可能表現更差，這表明在多芯片如MI300X/MI300A和B200/B300設備上編程時，可能需要新的微調方向。

然而，該算法的細節仍有進一步優化空間，以提高緩存命中率和主內存合并率。

致謝

特別感謝來自NUS團隊的覃含章教授，王昀鴻博士在MI100/MI250性能驗證中的合作，Zev Rekhter在MI300X性能驗證中的合作，范舒宜在H200驗證中的合作，以及BBuf在SGLang解決方案的討論和審閱。

請注意，這是SGLang社區的獨立工作。

作者介紹

本文作者王翼之前在Graphcore擔任機器學習專家，后加入美國知名半導體公司擔任AI架構師（SMTS主任工程師）。

參與貢獻諸多機器學習社區開源軟件，主要研究興趣在LLM訓練、推理的軟件棧優化，并應用計算機體系結構知識協同設計軟硬件解決方案。

參考資料：

[1]W. Fedus, B. Zoph, and N. Shazeer. Switch transformers: Scaling to trillion parameter models with simple and efficient sparsity. CoRR, abs/2101.03961, 2021. URL https://arxiv.org/abs/2101.03961.

[2]A. Q. Jiang, A. Sablayrolles, A. Mensch, C. Bamford, D. S. Chaplot, D. d. l. Casas, F. Bressand, G. Lengyel, G. Lample, L. Saulnier, et al. Mistral 7b. arXiv preprint arXiv:2310.06825, 2023.

[3]DeepSeek-AI. Deepseek-v2: A strong, economical, and efficient mixture-of-experts language model. CoRR, abs/2405.04434, 2024c. URL https://doi.org/10.48550/arXiv.2405.04434.

[4]DeepSeek V3 : https://arxiv.org/abs/2412.19437; Retrieved on 2025-03-18

[5]DeepSeek R1 : https://arxiv.org/pdf/2501.12948; Retrieved on 2025-03-18

[6]TransformerEngine : https://github.com/NVIDIA/TransformerEngine; Retrieved on 2025-03-18

[7]NV Group GEMM : https://github.com/yiakwy-xpu-ml-framework-team/NV_grouped_gemm; Retrieved on 2025-03-18

[8]FasterTransformer : https://github.com/NVIDIA/FasterTransformer; Retrieved on 2025-03-18

[9]CK Fused MoE V1 : ROCm/composable_kernel#1634

[10]AMD 3X MOE : https://rocm.blogs.amd.com/artificial-intelligence/DeepSeekR1-Part2/README.html

[11]Lean Wang and Huazuo Gao and Chenggang Zhao and Xu Sun and Damai Dai Auxiliary-Loss-Free Load Balancing Strategy for Mixture-of-Experts, 2024. URL https://arxiv.org/abs/2408.15664.

[12]PopART on chip TensorRemap : https://github.com/graphcore/popart/tree/sdk-release-3.4

[13] DeepSeek V3 Optimization based on AITER backend : sgl-project/sglang#4344

原文地址：

Github（https://shorturl.at/bEEn9），Hugging Face（https://shorturl.at/PMH3F）

特別聲明：以上內容(如有圖片或視頻亦包括在內)為自媒體平臺“網易號”用戶上傳并發布，本平臺僅提供信息存儲服務。

Notice: The content above (including the pictures and videos if any) is uploaded and posted by a user of NetEase Hao, which is a social media platform and only provides information storage services.