AMDGPU效能暴漲7倍，最佳化演算法首次開源！高效MoE支援任意專家數量

2025-05-20 03:42 新智元

新智元報道  

編輯：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團隊貢獻的。因此，它們採用了相同的方法進行部署。

CK

首個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.hppusing f_pipeline    = ck_tile::FusedMoeGemmPipeline_FlatmmUk<f_problem>;using f_partitioner = ck_tile::FusedMoeGemmTilePartitioner_Linear<f_shape>;using f_kernel      = ck_tile::FusedMoeGemmKernel<f_partitioner, f_pipeline, void>;const dim3 grids                       = f_kernel::GridSize(a);constexpr dim3 blocks                  = f_kernel::BlockSize();constexpr ck_tile::index_t kBlockPerCu = 1;staticint 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<blocks.x, kBlockPerCu>(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#L438CK_TILE_DEVICE voidmoe_align_block_size_kernel(...){constindex_t tid       = static_cast<index_t>(threadIdx.x);constindex_t start_idx = tid * tokens_per_thread;...#if 1if(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 arrangementpids_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_xcdstall_xcds = GRID_MN % NUM_XCDStall_xcds = NUM_XCDS if tall_xcds == 0 else tall_xcds# Compute current XCD and local pid within the XCDxcd = pid % NUM_XCDSlocal_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 xcdif xcd < tall_xcds:pid = xcd * pids_per_xcd + local_pidelse:pid = tall_xcds * pids_per_xcd + (xcd - tall_xcds) * (pids_per_xcd - 1) + local_pidif GROUP_SIZE_M == 1:pid_m = pid // num_pid_npid_n = pid % num_pid_nelse:num_pid_in_group = GROUP_SIZE_M * num_pid_ngroup_id = pid // num_pid_in_groupfirst_pid_m = group_id * GROUP_SIZE_Mgroup_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 clonehttps://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）


相關文章