Attention 백엔드

Attention 백엔드 (Attention Backend)

SGLang은 아주 다양한 attention 백엔드를 지원해요. 각각 장단점이 다르고, 필요에 따라 테스트해 볼 수 있어요.

최적의 attention 백엔드를 고르는 것은 성능을 극대화하는 데 핵심이에요. 백엔드마다 다양한 시나리오에서 뛰어나기 때문에, 모델·하드웨어·유스케이스에 따라 선택하세요. 모든 백엔드가 모든 플랫폼과 모델 아키텍처에서 지원되는 건 아니에요.

--attention-backend를 지정하지 않으면 SGLang은 하드웨어와 모델 아키텍처를 기반으로 가장 성능 좋은 백엔드를 자동 선택하려고 최선을 다해요.

지원 매트릭스

지원 매트릭스는 두 부분으로 나뉘어요: MHA(표준 attention)와 MLA(multi-head latent attention). MHA와 MLA의 핵심 차이에 대한 설명은 DeepSeek MLA 관련 SGLang 문서와 원본 DeepSeek MLA 논문을 참고하세요.

MHA 백엔드

백엔드	Page Size > 1 (네이티브)	FP8 KV 캐시	FP4 KV 캐시	Spec topk=1	Spec topk>1	슬라이딩 윈도우	멀티모달
FlashInfer	✅	✅	❌	✅	✅	✅	❌
FA3 (FlashAttention 3)	✅	✅	❌	✅	✅	✅	✅
FA4 (FlashAttention 4)	128	❌	✅	✅	✅	✅	✅
Triton	❌	✅	✅	✅	✅	✅	✅
Torch 네이티브 (SDPA)	❌	✅	✅	❌	❌	❌	✅
FlexAttention (PyTorch)	❌	❌	✅	❌	❌	❌	❌
TRTLLM MHA	16, 32 또는 64	✅	✅	✅	❌	✅	❌
Dual Chunk FlashAttention	✅	❌	❌	❌	❌	❌	❌
HPC-Ops	64	✅	❌	❌	❌	❌	❌
AITER (ROCm)	✅	✅	❌	✅	✅	✅	✅
Wave (ROCm)	✅	❌	❌	❌	❌	❌	❌
Ascend (NPU)	✅	❌	❌	✅	❌	✅	✅
Intel XPU	✅	❌	❌	❌	❌	✅	❌
Intel AMX (CPU)	❌	❌	❌	❌	❌	❌	❌

MLA 백엔드

백엔드	네이티브 페이지 크기	FP8 KV 캐시	FP4 KV 캐시	청크 프리픽스 캐시	Spec topk=1	Spec topk>1
FlashInfer MLA	1	❌	✅	✅	✅	❌
FlashMLA	64	✅	✅	✅	✅	❌
Cutlass MLA	128	✅	✅	✅	✅	❌
TRTLLM MLA (Blackwell)	32 또는 64	✅	✅	✅	✅	❌
CuteDSL MLA (Blackwell)	32 또는 64	✅	❌	✅	✅	❌
TokenSpeed MLA (Blackwell)	32 또는 64	✅ (필수)	❌	✅	✅	❌
FA3 (FlashAttention 3)	해당 없음	❌	❌	✅	✅	⚠️ (page_size=1만)
Triton	해당 없음	❌	❌	❌	✅	⚠️ (page_size=1만)
FA4	1	❌	✅	✅	❌	❌
Ascend MLA (NPU)	128	❌	❌	❌	❌	❌

멀티모달 attention은 --mm-attention-backend으로 선택돼요. "멀티모달" 열은 해당 백엔드 패밀리에 상응하는 멀티모달 구현이 존재하는지 나타내요.

DSA는 DeepSeek V3.2를 위해 특별히 설계됐어요. 자세한 내용은 DSA Attention 백엔드 섹션과 DeepSeek V3.2 배포 가이드를 참고하세요.

KV4 FA4 시나리오에서 FA4는 실행에 다른 --decode-attention-backend를 요구해요. trtllm_mha가 FA4와 호환되지 않는 것만 제외하면, 다른 decode 백엔드는 모두 표에 표시된 대로 동작해요.

추측 디코딩 topk: topk는 드래프트 모델에서 단계당 샘플링하는 드래프트 토큰 수예요. topk = 1은 고전 EAGLE을 따르고, topk > 1은 여러 브랜치를 탐색하며 드래프트·검증 경로 모두에서 백엔드 지원이 필요해요.

페이지 크기는 몇 개의 토큰을 하나의 KV 캐시 블록으로 묶을지 제어해요. 프리픽스 캐시가 효과를 보려면 토큰 수가 최소한 완전한 페이지 하나를 채워야 해요. 예를 들어 프롬프트가 32토큰이고 page_size = 64면 완전한 페이지를 채우지 못해 프리픽스 캐시에서 매칭될 수 없어요(페이지는 패딩 불가). 65토큰이고 page_size = 64면 64토큰의 첫 페이지만 캐시·매칭되고 나머지 1토큰은 버려져요. 최대 프리픽스 재사용(토큰 레벨 매칭)을 원하면 page_size = 1을 쓰세요. 다만 페이지 크기가 클수록 attention 커널 성능이 일반적으로 좋아지므로, 프리픽스 캐시 재사용이 중요하지 않다면 page_size > 1이 좋아요.

페이지에서 네이티브로 작동하지 않는 많은 백엔드는 페이지 테이블을 토큰당 인덱스로 확장해 래퍼 레이어에서 page_size > 1을 에뮬레이션할 수 있어요. "Page Size > 1 (네이티브)" 열은 진짜 인-커널 페이징을 나타내요. 일부 백엔드는 고정 네이티브 페이지 크기를 요구하고 다르게 축소/에뮬레이션할 수 없어요: TRTLLM MHA (16/32/64), TRTLLM MLA (32/64), CuteDSL MLA (32/64), FlashMLA (64), Cutlass MLA (128), Ascend (128), HPC-Ops (64).

MLA 페이지 크기 제약:

FlashInfer MLA: page_size = 1.
FlashMLA: page_size = 64.
Cutlass MLA: page_size = 128.
TRTLLM MLA: page_size ∈ {32, 64}.
CuteDSL MLA: page_size ∈ {32, 64} (decode 전용; 미설정 시 prefill은 trtllm_mla로 폴백).
TokenSpeed MLA: page_size ∈ {32, 64} (Blackwell SM100/SM12x 전용; --kv-cache-dtype fp8_e4m3 필요).

GDN Attention 백엔드

GDN(Gated Delta Network)은 O(n) 복잡도의 선형 attention 메커니즘으로, GDN 선형 attention 레이어와 표준 full attention 레이어를 번갈아 쓰는 하이브리드 모델에서 사용돼요. GDN은 --attention-backend로 선택되지 않고, 모델 아키텍처가 요구할 때 자동으로 활성화돼요 (예: Qwen 3.5, Qwen 3 Next, Jet Nemotron, Jet VLM).

GDN 선형 attention 레이어는 자체 커널 백엔드를 갖고 있으며, --linear-attn-backend(기본: triton)로 선택돼요. 단계를 기준으로 --linear-attn-decode-backend와 --linear-attn-prefill-backend로 커널을 재정의할 수 있어요.

SM100/SM103과 CUDA 13+에서, per-phase 오버라이드가 미설정이고 기본 선형-attention 백엔드가 Triton이며, 순환 상태가 BF16, key/value 헤드 차원이 128이고, 동적 청킹과 page-major KV 레이아웃이 비활성이며, --chunked-prefill-size가 1에서 8192 사이일 때 SGLang은 GDN prefill에 FlashInfer를 자동 선택해요. Radix 캐싱은 비활성화되거나 no_buffer, extra_buffer, extra_buffer_lazy를 사용할 수 있어요(extra-buffer 경로는 상태 체크포인트를 사용해요).

백엔드	Decode	Prefill / Extend	Spec Decoding (Target Verify)
Triton (CUDA)	✅	✅	✅
Triton (AMD/ROCm)	✅	✅	✅
Triton (NPU)	✅	✅	❌
Triton (CPU)	✅	✅	❌
CuTe DSL (CUDA 전용)	✅	❌	❌
FlashInfer (CUDA, SM90/SM100/SM103)	✅	✅	✅ linear chain; 트리는 Triton으로 폴백

GDN 모델은 하이브리드입니다: full-attention 레이어는 여전히 표준 --attention-backend가 필요해요. 하이브리드 GDN 모델에서 full-attention 백엔드의 플랫폼 제약:

Blackwell SM120 (예: RTX PRO 6000 Blackwell): prefill/full attention엔 triton 또는 flashinfer; trtllm_mha는 --decode-attention-backend로만 지원.

다른 Blackwell 변형 (SM100 B200/GB200 포함): triton, trtllm_mha, 또는 fa4만.

NPU (Ascend): ascend만.

AMD (ROCm): triton 권장.

기타 CUDA (Hopper, Ampere 등): 자동 선택 동작; 특별한 제약 없음.

DSA Attention 백엔드

DSA(DeepSeek Sparse Attention)는 DeepSeek V3.2가 사용하는 네이티브 희소 attention 메커니즘이에요. 모델 아키텍처가 요구할 때 자동으로 활성화되고 --attention-backend dsa(비권장 별칭: nsa)로 선택돼요.

내부적으로 DSA 백엔드는 prefill과 decode 단계에 서로 다른 하위 백엔드로 디스패치해요. --dsa-prefill-backend와 --dsa-decode-backend로 오버라이드할 수 있어요:

하위 백엔드	Prefill	Decode	참고
flashmla_sparse	✅	✅	Hopper·Blackwell(BF16)에서 기본 prefill
flashmla_sparse_q8	✅	❌	Hopper(SM90)에서 네이티브 FP8(q8×kv8) 희소 prefill; `--kv-cache-dtype fp8_e4m3` 필요
flashmla_kv	✅	✅	Hopper에서 FP8 기본(prefill + decode)
flashmla_auto	✅	❌	KV 캐시 dtype으로 flashmla_sparse 또는 flashmla_kv 선택
fa3	✅	✅	Hopper(BF16)에서 기본 decode
trtllm	✅	✅	Blackwell(BF16)에서 기본 decode; Blackwell에서 FP8 기본(prefill + decode)
tilelang	✅	✅	AMD(ROCm)에서 기본
aiter	✅	✅	AMD 특화 커널 라이브러리 (aiter 패키지 필요)

`aiter`(ROCm)에서의 Gemma 4

Gemma 4의 하이브리드 attention은 슬라이딩 윈도우 레이어(head_dim 256)와 full-attention 레이어(head_dim 512)를 섞고 레이어 간 KV를 공유해요. aiter 백엔드는 이를 지원하며 Gemma 4용으로 허용 목록에 있어, MI300X/MI350X에서 --attention-backend aiter로 서빙할 수 있어요. 내부적으로 aiter는 head_dim-512 prefill을 Triton unified_attention 커널로 라우팅하고(CK 배치-prefill 경로는 head_dim을 256으로 제한), 각 레이어의 softmax 스케일을 존중하며, speculative-decode 드래프트 워커의 슬라이딩 윈도우 KV 풀을 해결해요. 비-슬라이딩 윈도우 모델(head_dim ≤ 256)은 영향받지 않고 CK 경로를 계속 사용해요.

Gemma 4를 aiter에서 쓸 때는 bf16 KV 캐시(--kv-cache-dtype auto)를 사용하세요. --kv-cache-dtype fp8_e4m3는 쓰지 마세요: 이 경로에서 fp8 KV를 쓰면 배치 decode가 현재 잘못된 출력을 반환해요(기존 aiter fp8-decode 이슈, 별도로 추적됨). bf16 KV는 모든 배치 크기에서 정확해요.

배포 예시는 DeepSeek V3.2 배포 가이드를 참고하세요.

하이브리드 attention (prefill·decode용 다른 백엔드) (실험)

하이브리드 attention은 실험 기능이에요.

prefill과 decode에 서로 다른 attention 백엔드를 섞어 쓸 수 있어요. 한 백엔드가 prefill에, 다른 백엔드가 decode에 뛰어날 때 유용해요. 구현 세부사항은 python/sglang/srt/layers/attention/hybrid_attn_backend.py를 참고하세요.

# 예시: FA4로 Prefill, TRTLLM MLA로 Decode (Blackwell)
python3 -m sglang.launch_server \
  --model-path nvidia/DeepSeek-R1-FP4 \
  --tp 8 \
  --attention-backend trtllm_mla \
  --moe-runner-backend flashinfer_trtllm \
  --quantization modelopt_fp4 \
  --prefill-attention-backend fa4

하이브리드 attention과 추측 디코딩

하이브리드 attention은 추측 디코딩과도 함께 작동해요. 드래프트 디코딩과 타깃 검증에 사용되는 백엔드는 --speculative-attention-mode에 따라 달라져요:

--speculative-attention-mode decode (권장): draft/verify는 decode 백엔드 사용.
--speculative-attention-mode prefill (기본): draft/verify는 prefill 백엔드 사용.

하이브리드 attention을 추측 디코딩과 결합할 때의 제약:

어떤 attention 백엔드든 trtllm_mha면 추측 디코딩은 --speculative-eagle-topk 1만 지원해요.
--page-size > 1 + --speculative-eagle-topk > 1인 paged MHA 백엔드에선 flashinfer만 지원돼요.
CUDA 그래프: decode 백엔드는 항상 캡처되고, prefill 백엔드는 --speculative-attention-mode prefill일 때만 캡처돼요.

--prefill-attention-backend 또는 --decode-attention-backend 중 하나만 설정하면, 미지정 단계는 --attention-backend를 상속해요. 둘 다 지정되고 다르면 SGLang은 단계별로 선택한 백엔드로 디스패치하는 하이브리드 래퍼를 자동 활성화해요.