[sglang] SM120 Blackwell에서 DeepSeek-V4 모델 서빙 최적화: FlashInfer MXFP4 MoE 도입 및 메모리 절감
PR 링크: sgl-project/sglang#30272 상태: Merged | 변경: +506 / -237
들어가며
최근 대규모 언어 모델(LLM)의 발전과 함께, 다양한 하드웨어 아키텍처에서의 효율적인 서빙은 매우 중요한 과제가 되었습니다. 특히 NVIDIA의 최신 GPU 아키텍처인 Blackwell(SM120) 환경에서 DeepSeek-V4와 같은 복잡한 Mixture-of-Experts (MoE) 모델을 서빙하는 것은 여러 기술적 난관에 부딪혔습니다. 본 PR은 이러한 문제를 해결하고, SM120 GPU에서 DeepSeek-V4 모델의 서빙 성능과 안정성을 획기적으로 개선하는 것을 목표로 합니다.
이 PR은 특히 다음과 같은 문제점들을 해결합니다:
- MoE Forward Crash: Marlin MoE 백엔드에서 EP(Expert Parallelism) 환경에서 마스크된
-1expert 블록을 잘못 처리하여 발생하는 크래시. - KV Cache OOM: FP32 로더 컨테이너가 expert 스케일을 과도하게 확장하여 약 16GB/rank의 메모리를 추가로 소모하는 문제.
- Sparse-Prefill Kernel 미지원: SM120 아키텍처가
sparse-prefill커널을 지원하지 않아 발생하는 오류. - Prefill OOM: SM120 래퍼가 임의의 extend 배치에 대해 split-K 디코드 커널을 사용하여 토큰당 약 1MB의 scratch 공간을 할당하는 문제.
- Indexer OOM: Torch fallback이 FP32로 gather된 index KV를 전체 청크에 대해 생성하여 발생하는 메모리 부족 문제.
- 비효율적인 MoE 경로: DSv4 expert와 FlashInfer SM120 커널이 MXFP4 × MXFP8을 네이티브로 지원함에도 불구하고, Marlin은 W4A16 경로를 사용하는 비효율성.
이 PR은 이러한 문제들을 해결하기 위해 FlashInfer MXFP4 MoE 백엔드를 SM120 기본값으로 설정하고, Marlin MoE 경로의 메모리 사용량을 최적화하며, Attention 및 Indexer 관련 커널들을 개선합니다.
코드 분석
1. FlashInfer MXFP4 MoE (SM120 기본값)
이 PR의 핵심 변경 사항 중 하나는 SM120 GPU에서 DeepSeek-V4 모델 서빙 시 기본 MoE 백엔드를 FlashInfer MXFP4로 설정하는 것입니다. 이는 기존의 Marlin W4A16 기반 경로보다 훨씬 효율적입니다.
변경 사항:
- DSv4 어댑터 추가: FlashInfer CUTLASS MoE를 위한 DSv4 어댑터가 추가되었습니다. SM120에서는 라우팅된 expert 가중치를 MXFP4로 유지하고, 스케일 메타데이터만 E8M0으로 인터리빙하며, 활성화(activation)를 동적으로 MXFP8으로 양자화합니다. 이는 W4A16 또는 디퀀타이즈된 expert 가중치 복사본을 생성하지 않아 메모리를 절약합니다.
- FlashInfer 러너 확장: MXFP8 × MXFP4 호출 계약(activation scales, global weight scale, SwiGLU clamp, normalized top-k tensors 포함)을 지원하도록 확장되었습니다.
- 기본값 설정:
--moe-runner-backend auto옵션이 SM120의 DeepSeek-V4에 대해flashinfer_mxfp4로 라우팅됩니다. 명시적인 백엔드 선택은 여전히 우선권을 가집니다. - SM90 호환성 유지: SM90용 FlashInfer 혼합 입력 W4A16 경로를 유지하고, SM90/SM120 어댑터 커버리지를 추가했습니다.
- FlashInfer 버전 업그레이드:
flashinfer_python,flashinfer_cubin, Docker JIT-cache 버전이 0.6.12에서 0.6.14로 업그레이드되어 SM120 MXFP8 활성화 스케일링 및 sparse-MLA API를 사용할 수 있게 되었습니다.
코드 인용:
python/sglang/srt/arg_groups/overrides.py 파일에서 _deepseek_v4_sm120_moe 함수가 수정되었습니다:
--- a/python/sglang/srt/arg_groups/overrides.py
+++ b/python/sglang/srt/arg_groups/overrides.py
@@ -1434,16 +1434,15 @@
@register_post_process
def _deepseek_v4_sm120_moe(view: Any) -> dict:
- """Slot pass in the DeepSeek V4 validation branch: SM120 lacks
- tcgen05/TMEM, fall back to the marlin MoE runner (reads the
- mid-resolution moe_runner_backend, after the dispatch-time nvfp4
- default)."""
+ """Default DeepSeek V4 MXFP4 experts to FlashInfer CUTLASS on SM120."""
hf_config = view.get_model_config().hf_config
if hf_config.architectures[0] != "DeepseekV4ForCausalLM":
return {}
if is_sm120_supported() and view.moe_runner_backend == "auto":
- logger.info("Use marlin as MoE runner backend on SM120 for DeepseekV4")
- return {"moe_runner_backend": "marlin"}
+ logger.info(
+ "Use flashinfer_mxfp4 as MoE runner backend on SM120 for DeepseekV4"
+ )
+ return {"moe_runner_backend": "flashinfer_mxfp4"}
return {}
이 변경으로 인해 SM120 환경에서 --moe-runner-backend auto 설정 시 Marlin 대신 flashinfer_mxfp4가 기본으로 사용됩니다.
2. Marlin MoE 정확성 및 메모리 최적화
FlashInfer로 전환하면서도, 기존 Marlin MoE 경로의 버그 수정 및 메모리 최적화는 계속 진행되었습니다. 이는 다른 모델이나 환경에서 Marlin을 사용할 때의 안정성과 효율성을 높입니다.
변경 사항:
marlin_template.h:expert_id == -1인 블록을is_ep플래그와 관계없이 무조건 건너뛰도록 수정하여 EP 환경에서의 크래시를 방지합니다.fused_marlin_moe.py: 건너뛴 행(row)이 쓰여지지 않은 상태로 남아 가중 합계에 소비되는 문제를 해결하기 위해, 별칭(aliased) 중간 캐시를 0으로 초기화합니다.marlin_utils_fp4.py: Marlin 준비 후 로더 형식의 스케일 파라미터를 해제하여 메모리를 절약합니다.mxfp4_marlin_moe.py: 로더 스케일을 FP32로 확장하는 대신, 체크포인트의 E8M0 dtype으로 등록하여 메모리 사용량을 줄입니다. 리뷰어ormandj의 의견에 따르면 이 변경으로 인해 로드 시 메모리가 약 16GB/rank 절감되었습니다.
코드 인용:
python/sglang/jit_kernel/csrc/gemm/marlin_moe/marlin_template.h의 수정:
--- a/python/sglang/jit_kernel/csrc/gemm/marlin_moe/marlin_template.h
+++ b/python/sglang/jit_kernel/csrc/gemm/marlin_moe/marlin_template.h
@@ -375,14 +375,11 @@ __global__ void Marlin(
is_zp_float ? prob_n * prob_k / group_size / 8 : prob_n * prob_k / group_size / (pack_factor * 4);
const int b_bias_expert_stride = prob_n / 8;
- // parallel: num valid moe blocks
int num_tokens_past_padded = num_tokens_past_padded_ptr[0];
int parallel = num_tokens_past_padded / moe_block_size;
int num_valid_blocks = parallel;
- if (is_ep) {
- for (int i = 0; i < parallel; i++) {
- if (expert_ids_ptr[i] == -1) num_valid_blocks--;
- }
+ for (int i = 0; i < parallel; i++) {
+ if (expert_ids_ptr[i] == -1) num_valid_blocks--;
}
int num_invalid_blocks = parallel - num_valid_blocks;
parallel = num_valid_blocks;
3. Attention 최적화
SM120에서 sparse-prefill 커널이 지원되지 않는 문제를 해결하고, extend 배치를 FlashInfer sparse MLA를 통해 처리하도록 변경했습니다.
변경 사항:
deepseek_v4_backend.py: SM120에서 지원되지 않는sparse-prefill경로를 비활성화하고, extend 배치를 FlashInfer sparse MLA로 라우팅합니다.flash_mla_sm120.py: FlashInfer의 디스패처 진입점을 호출합니다. 디코드 크기 배치는 split-K를 사용하고, 더 큰 extend 배치는 토큰당 split-K scratch 공간 없이 일반적인 paged-attention 커널을 사용합니다. 이는 Prefill OOM 문제를 해결합니다.
코드 인용:
python/sglang/srt/layers/attention/deepseek_v4_backend.py의 수정:
--- a/python/sglang/srt/layers/attention/deepseek_v4_backend.py
+++ b/python/sglang/srt/layers/attention/deepseek_v4_backend.py
@@ -1688,9 +1688,14 @@ def match_num_queries(x, value):
extra_indices.shape[-1] % 64 == 0
), f"{extra_indices.shape=}'s last dimension is not aligned to 64"
- if forward_batch.forward_mode.is_extend_without_speculative() and (
- q.shape[0] > _LARGE_INDEXER_QUERY_THRESHOLD
- or envs.SGLANG_OPT_FLASHMLA_SPARSE_PREFILL.get()
+ # sparse_prefill_fwd does not support SM120.
+ if (
+ forward_batch.forward_mode.is_extend_without_speculative()
+ and not _is_sm120
+ and (
+ q.shape[0] > _LARGE_INDEXER_QUERY_THRESHOLD
+ or envs.SGLANG_OPT_FLASHMLA_SPARSE_PREFILL.get()
+ )
):
return self._forward_prefill_sparse(
q=q,
python/sglang/kernels/ops/attention/flash_mla_sm120.py의 수정:
--- a/python/sglang/kernels/ops/attention/flash_mla_sm120.py
+++ b/python/sglang/kernels/ops/attention/flash_mla_sm120.py
@@ -400,14 +400,19 @@ def _flash_mla_flashinfer(
extra_indices,
extra_topk_length,
):
- """FlashInfer SM120 sparse MLA via sparse_mla_sm120_decode_dsv4.
+
+ """FlashInfer SM120 sparse MLA via the paged-attention dispatcher.
SGLang SWA pool uses page_size=256 (footer format: 256*576 bytes data + 256*8 bytes scale).
FlashInfer decode_dsv4 fast path requires page_block_size=64 (footer: 64*576 + 64*8).
We split 256-token pages into 4 virtual 64-token pages.
Token indices are invariant under page-split (identity mapping).
"""
- from flashinfer.mla._sparse_mla_sm120 import sparse_mla_sm120_decode_dsv4
+ from flashinfer.mla._sparse_mla_sm120 import (
+ _DECODE_MAX_TOKENS as _FI_DECODE_MAX_TOKENS,
+ )
+ from flashinfer.mla._sparse_mla_sm120 import (
+ _sparse_mla_sm120_paged_attention,
+ )
B, _, H, D = q.shape # (batch, 1, num_heads, head_dim)
dev = q.device
@@ -435,32 +440,37 @@ def _flash_mla_flashinfer(
output = torch.empty(B, H, head_dim_v, dtype=torch.bfloat16, device=dev)
out_lse = torch.empty(B, H, dtype=torch.float32, device=dev)
- # Pre-allocate split-K scratch for decode-dsv4 fast path.
- topk = idx.shape[-1]
- extra_topk = extra_idx.shape[-1] if extra_idx is not None else 0
- _BI = 64
- num_splits = (topk + _BI - 1) // _BI + (
- (extra_topk + _BI - 1) // _BI if extra_topk > 0 else 0
- )
- mid_out = torch.empty(
- B, H, num_splits, head_dim_v, dtype=torch.bfloat16, device=dev
- )
- mid_lse = torch.empty(B, H, num_splits, dtype=torch.float32, device=dev)
-
- sparse_mla_sm120_decode_dsv4(
- q=q.squeeze(1) if q.ndim == 4 else q,
- kv_cache=kv_64,
- indices=idx,
- mid_out=mid_out,
- mid_lse=mid_lse,
- output=output,
- out_lse=out_lse,
- sm_scale=softmax_scale,
+ # Use split-K for decode-sized batches and paged attention otherwise.
+ if B <= _FI_DECODE_MAX_TOKENS:
+ topk = idx.shape[-1]
+ extra_topk = extra_idx.shape[-1] if extra_idx is not None else 0
+ _BI = 64
+ num_splits = (topk + _BI - 1) // _BI + (
+ (extra_topk + _BI - 1) // _BI if extra_topk > 0 else 0
+ )
+ mid_out = torch.empty(
+ B, H, num_splits, head_dim_v, dtype=torch.bfloat16, device=dev
+ )
+ mid_lse = torch.empty(B, H, num_splits, dtype=torch.float32, device=dev)
+ else:
+ mid_out = None
+ mid_lse = None
+
+ _sparse_mla_sm120_paged_attention(
+ q.squeeze(1) if q.ndim == 4 else q,
+ kv_64,
+ idx,
+ output,
+ out_lse,
+ softmax_scale,
+ d_v=head_dim_v,
topk_length=topk_length,
attn_sink=attn_sink,
extra_kv_cache=extra_kv_64,
extra_indices=extra_idx,
extra_topk_length=extra_topk_length,
+ mid_out=mid_out,
+ mid_lse=mid_lse,
)
return (output.unsqueeze(1), None)
4. Indexer 최적화
Indexer 부분에서도 메모리 사용량과 성능을 개선하기 위한 수정이 이루어졌습니다. 특히 Torch fallback 시 메모리 사용량을 제한하고, TileLang 커널의 버퍼 구성을 변경했습니다.
변경 사항:
tilelang_kernel.py:fp8_paged_mqa_logits스테이징 버퍼를 2차원으로 만들어 TileLang의 GEMM 레이아웃 하향 조절(lowering)이 행 크기(row size)를 결정할 수 있도록 합니다.server_args.py: SM120에서 TileLang 인덱서를 기본값으로 활성화하고 Torch fallback은 유지합니다.indexer.py: Torch fallback의 메모리 사용량을 제한하기 위해 참조 쿼리 차원(reference query dimension)을 1024 토큰으로 청크(chunk)합니다.
코드 인용:
python/sglang/kernels/ops/attention/dsa/tilelang_kernel.py의 수정:
--- a/python/sglang/kernels/ops/attention/dsa/tilelang_kernel.py
+++ b/python/sglang/kernels/ops/attention/dsa/tilelang_kernel.py
@@ -1443,11 +1443,14 @@ def fp8_paged_mqa_logits(
for j in T.Pipelined(n_iters, num_stages=2):
i = i_start + j
page = page_table[bx, i]
- k_smem_u8 = T.alloc_shared((B * D,), UINT8)
- T.copy(kvcache_u8[page, 0:SCALE_OFFSET], k_smem_u8)
+ k_smem_u8 = T.alloc_shared((1, B * D), UINT8)
+ T.copy(kvcache_u8[page : page + 1, 0:SCALE_OFFSET], k_smem_u8)
k_smem = T.view(k_smem_u8, (B, D), FP8)
- k_s_smem_u8 = T.alloc_shared((B * 4,), UINT8)
- T.copy(kvcache_u8[page, SCALE_OFFSET:BLOCK_BYTES], k_s_smem_u8)
+ k_s_smem_u8 = T.alloc_shared((1, B * 4), UINT8)
+ T.copy(
+ kvcache_u8[page : page + 1, SCALE_OFFSET:BLOCK_BYTES],
+ k_s_smem_u8,
+ )
k_s_smem = T.view(k_s_smem_u8, (B,), FP32)
k_s_frag = T.alloc_fragment((B,), FP32)
T.copy(k_s_smem, k_s_frag)
python/sglang/srt/layers/attention/dsv4/indexer.py의 수정:
--- a/python/sglang/srt/layers/attention/dsv4/indexer.py
+++ b/python/sglang/srt/layers/attention/dsv4/indexer.py
@@ -174,6 +174,25 @@ def fp8_paged_mqa_logits_torch_sm120(
block_size = kvcache_fp8.shape[1]
device = q_fp8.device
+ _QUERY_CHUNK = 1024
+ if batch_size > _QUERY_CHUNK:
+ return torch.cat(
+ [
+ fp8_paged_mqa_logits_torch_sm120(
+ q_fp8[start : start + _QUERY_CHUNK],
+ kvcache_fp8,
+ weight[start : start + _QUERY_CHUNK],
+ kv_scale[start : start + _QUERY_CHUNK],
+ kv_offset[start : start + _QUERY_CHUNK],
+ block_size,
+ device,
+ )
+ for start in range(_QUERY_CHUNK, batch_size, _QUERY_CHUNK)
+ ]
+ )
+
_QUERY_CHUNK = 1024
if batch_size > _QUERY_CHUNK:
return torch.cat(
5. 문서 업데이트
SM120 환경에서의 DeepSeek-V4 서빙 설정 및 권장 사항이 문서에 반영되었습니다.
코드 인용:
docs_new/cookbook/autoregressive/DeepSeek/DeepSeek-V4.mdx의 수정:
--- a/docs_new/cookbook/autoregressive/DeepSeek/DeepSeek-V4.mdx
+++ b/docs_new/cookbook/autoregressive/DeepSeek/DeepSeek-V4.mdx
@@ -149,7 +149,7 @@ import { Playground } from "/src/snippets/_playground.jsx";
<td style={{padding: "9px 12px", fontWeight: 500, backgroundColor: "rgba(255,255,255,0.02)"}}><strong><a href="https://huggingface.co/deepseek-ai/DeepSeek-V4-Flash">DeepSeek-V4-Flash</a></strong></td>
<td style={{padding: "9px 12px", textAlign: "right", backgroundColor: "rgba(255,255,255,0.05)"}}><strong>284B</strong></td>
<td style={{padding: "9px 12px", textAlign: "right", backgroundColor: "rgba(255,255,255,0.02)"}}>13B</td>
- <td style={{padding: "9px 12px", backgroundColor: "rgba(255,255,255,0.05)"}}>single-node serving on B200 / B300 / GB200 / GB300 / H200 (TP=4); H100 (TP=8)</td>
+ <td style={{padding: "9px 12px", backgroundColor: "rgba(255,255,255,0.05)"}}>single-node serving on B200 / B300 / GB200 / GB300 / H200 (TP=4); RTX PRO 6000 (TP=2); H100 (TP=8)</td>
</tr>
<tr>
<td style={{padding: "9px 12px", fontWeight: 500, backgroundColor: "rgba(255,255,255,0.02)"}}><strong><a href="https://huggingface.co/deepseek-ai/DeepSeek-V4-Pro">DeepSeek-V4-Pro</a></strong></td>
@@ -293,9 +293,8 @@ TCP, which can lead to garbled KV transfer on large checkpoints.
**RTX PRO 6000 (SM120 / Blackwell Desktop) note**
-RTX PRO 6000 (96 GB) runs **Flash only** — V4-Pro doesn't fit on 8× 96 GB. It uses the
-**low-latency / TP-only** recipe (TP=4, single node) with the **Marlin** W4A16 MoE runner and
-`--mem-fraction-static 0.70`; the Deploy panel greys out the other recipes for this card.
+RTX PRO 6000 (96 GB) runs **Flash only** with the FlashInfer MXFP4 MoE runner.
V4-Pro doesn't fit on 8× 96 GB; the Deploy panel greys out unsupported recipes.
HiCache and MegaMoE are **not** supported on RTX PRO 6000.
**AMD (MI300X / MI355X) note**
docs_new/src/snippets/configs/deepseek-ai/deepseek-v4.jsx의 수정:
--- a/docs_new/src/snippets/configs/deepseek-ai/deepseek-v4.jsx
+++ b/docs_new/src/snippets/configs/deepseek-ai/deepseek-v4.jsx
@@ -1421,7 +1421,6 @@ sgl-eval run aime25 \
// ====================================================================
// RTX PRO 6000 (SM120 / Blackwell Desktop) — Flash + low-latency only
- // (V4-Pro doesn't fit on 8× 96 GB); TP-only, Marlin MoE runner.
// ====================================================================
{
match: { hw: "rtx6000", variant: "flash", quant: "fp4", strategy: "low-latency", nodes: "single" },
@@ -1430,9 +1429,9 @@ sgl-eval run aime25 \
flags: [
"--trust-remote-code",
"--model-path {{MODEL_NAME}}",
- "--tp 4",
- "--moe-runner-backend marlin",
- "--mem-fraction-static 0.70",
+ "--tp 2",
+ "--moe-runner-backend flashinfer_mxfp4",
+ "--mem-fraction-static 0.92",
"--cuda-graph-max-bs-decode 32",
"--host {{HOST_IP}}",
"--port {{PORT}}",
왜 이게 좋은가
이 PR은 여러 측면에서 상당한 개선을 가져왔습니다.
성능 향상
PR 설명에 포함된 속도 테스트 결과에 따르면, SM120 환경(2x RTX PRO 6000)에서 DeepSeek-V4-Flash TP2/EP2 모델 서빙 시 FlashInfer MXFP4 MoE 백엔드를 사용했을 때, 기존 Marlin Inferlab 실행 대비 다음과 같은 성능 향상을 보였습니다:
| concurrency | FlashInfer output tok/s | Marlin output tok/s | change |
|---|---|---|---|
| 1 | 53.9 | 50.9 | +5.8% |
| 2 | 97.8 | 92.8 | +5.5% |
| 4 | 167.1 | 156.7 | +6.7% |
| 8 | 255.7 | 236.8 | +8.0% |
| 16 | 306.5 | 280.2 | +9.4% |
| 32 | 309.5 | 283.3 | +9.3% |
또한, TTFT(Time To First Token)에서도 개선이 관찰되었습니다:
| concurrency | FlashInfer mean TTFT | Marlin mean TTFT | change |
|---|---|---|---|
| 1 | 1.00 s | 1.17 s | -14.4% |
| 2 | 1.79 s | 2.10 s | -15.1% |
| 4 | 2.91 s | 3.42 s | -15.0% |
| 8 | 4.91 s | 5.78 s | -15.1% |
| 16 | 10.47 s | 12.07 s | -13.3% |
| 32 | 52.13 s | 57.47 s | -9.3% |
이 결과는 FlashInfer MXFP4 MoE 경로가 SM120 아키텍처에 더 최적화되어 있으며, 특히 고부하(high concurrency) 환경에서 더 나은 처리량과 응답 시간을 제공함을 보여줍니다.
메모리 절감
이 PR은 메모리 사용량도 크게 줄였습니다. 특히 E8M0 로더 컨테이너 수정으로 인해 TP2/EP2 환경에서의 모델 로드 시 메모리 사용량이 91.99 GB/rank에서 75.06 GB/rank로 약 17 GB/rank 감소했습니다. 이는 KV 캐시 풀을 위한 더 많은 헤드룸을 확보하여, 더 긴 시퀀스나 더 많은 동시 요청을 처리할 수 있게 합니다.
리뷰어 ormandj의 분석에 따르면, Marlin 관련 메모리 최적화(예: 로더 스케일 파라미터 해제, E8M0 dtype 등록)는 약 16GB + 12GB = 28GB/rank의 절감 효과를 가져올 수 있습니다. 이는 MoE 모델의 높은 메모리 요구량을 고려할 때 매우 중요한 개선입니다.
안정성 향상
Marlin MoE의 -1 expert 블록 처리 오류 수정, Torch fallback 청킹을 통한 OOM 방지 등은 모델 서빙의 안정성을 크게 향상시킵니다. 특히 EP 환경에서의 크래시 방지는 프로덕션 환경에서 매우 중요합니다.
일반적 교훈
- 하드웨어 특화 최적화의 중요성: 최신 GPU 아키텍처(SM120)의 특성을 최대한 활용하기 위해 네이티브 지원되는 데이터 타입(MXFP4, MXFP8)과 커널(FlashInfer)을 적극적으로 도입하는 것이 성능 향상의 핵심입니다.
- 메모리 최적화는 성능만큼 중요: LLM, 특히 MoE 모델은 메모리 집약적입니다. KV 캐시 및 모델 가중치 로딩 시 메모리 사용량을 줄이는 것은 서빙 용량과 효율성을 직접적으로 향상시킵니다.
- 버그 수정과 성능 개선의 병행: 새로운 기능을 도입하면서 기존 코드의 버그를 수정하고 메모리 사용량을 최적화하는 것은 전체 시스템의 안정성과 성능을 동시에 높이는 효과적인 전략입니다.
- 다양한 백엔드 지원 및 자동 선택:
auto옵션을 통해 하드웨어에 맞는 최적의 백엔드를 자동으로 선택하도록 하는 것은 사용자 편의성과 성능을 모두 잡는 좋은 방법입니다. 필요에 따라 명시적 선택도 가능해야 합니다.
References
- FlashInfer Python API
- Torch.cat
- NVIDIA Blackwell Architecture
- DeepSeek-V4-Flash
- Marlin MoE (참고용, 본 PR에서 직접 수정된 코드는 sglang 레포에 있음)
참고 자료
- https://github.com/Dao-AILab/flash-inference/blob/main/python/flashinfer/mla/_sparse_mla_sm120.py
- https://pytorch.org/docs/stable/generated/torch.cat.html
- https://www.nvidia.com/en-us/data-center/technologies/blackwell-gpu-architecture/
- https://huggingface.co/deepseek-ai/DeepSeek-V4-Flash
⚠️ 알림: 이 분석은 AI가 실제 코드 diff를 기반으로 작성했습니다.
관련 포스트
PR Analysis 의 다른글
- 이전글 [onnxruntime] ONNX Runtime: Arm64 KleidiAI 기반 FP16 GEMM 및 Convolution 최적화
- 현재글 : [sglang] SM120 Blackwell에서 DeepSeek-V4 모델 서빙 최적화: FlashInfer MXFP4 MoE 도입 및 메모리 절감
- 다음글 [sglang] SGLang DFlash 최적화: 호스트-디바이스 동기화 제거를 통한 추론 성능 향상
댓글