Q: FreeToken으로 서빙할 때 Qwen3.5/Qwen3.6 RMSNorm 가중치의 bf16 정밀도 손실을 해결하는 방법은?

버그: FreeToken의 `models/qwen3_5_moe` 로더는 로드 시점에 Gemma 스타일 `+1`을 bf16 체크포인트 텐서에 합산하여(`tensor + 1.0`) `bf16(1 + w)`를 저장합니다. 반면 Transformers는 제로 중심 가중치를 유지하고 `(1 + w)`를 fp32로 계산하므로, 이 합산 과정에서 작은 가중치가 반올림으로 소실됩니다. `nvidia/Qwen3.6-35B-A3B-NVFP4`에서는 171,008개의 norm 승수 중 77.4%가 변경되며, 평균 상대 오차는 1.4e-3, 최대는 3.9e-3(bf16 ulp의 절반)으로, 모든 토큰에서 `input_layernorm`, `post_attention_layernorm`, `self_attn.q_norm`, `self_attn.k_norm` 및 최종 `model.norm`에 영향을 미칩니다. 수정 방법: 원래의 제로 중심 bf16 가중치 `w`를 그대로 저장하고(로더에서 1을 더하지 마세요) 런타임에 fp32로 `+1`을 적용하세요. 기존 `GemmaPlusOneRMSNorm` 레이어(`layers/norm.py`)를 사용하거나 `sgl_kernel.rmsnorm(x, bf16(1+w))` 대신 `sgl_kernel.gemma_rmsnorm(x, w)`를 호출하면 됩니다. 이는 `qwen4_exp` 로더에서 이미 사용 중인 방식과 일치하며, fp64 레퍼런스와의 bf16 일치율을 99.999%로 복원합니다(잔여 오차는 단순한 출력 반올림이며, rel L2는 이전 3.0e-3에서 ~1.7e-3으로 감소).

Question 1

FreeToken이 --allowed-local-media-path 허용 목록에서 'file://localhost' 이미지 URI를 거부하는 이유는 무엇인가요?

Accepted Answer

이것은 `_load_local_media`(`python/freetoken/mm/media.py`)의 URI 파싱 버그입니다. 이 함수는 URI의 `netloc`과 `path`를 연결하여 파일 경로를 생성하므로, `file://localhost/path/image.png`가 `/path/image.png` 대신 상대 경로 `localhost/path/image.png`가 되며, 이로 인해 resolved-path allowlist 검사에 실패하여 `ValueError`가 발생합니다. 수정 방법: RFC 8089에 따라 비어 있는 authority와 `localhost` authority를 (대소문자 구분 없이) 로컬로 처리하고, 그 밖의 모든 비로컬 authority는 거부하며, `url2pathname`으로 경로를 디코딩하고, 기존의 `Path.resolve()` strict-root allowlist 검사는 그대로 유지하세요. 패치될 때까지의 우회 방법: `file://localhost/...` 대신 정규 형태인 `file:///path/image.png`(빈 authority)를 사용하세요.

Question 2

왜 FreeToken이 RTX 5090 / 5070 Ti (sm_120)에서 첫 번째 TMA 커널 실행 후 OOM이 발생하거나 1.3-3.3 GiB의 VRAM을 잃는가?

Accepted Answer

sm_120 GPU에서 Triton 3.8 미만 버전은 TMA 로드(TensorDescriptor)를 첫 번째 커널 실행 시 모든 상주 스레드에 대해 스레드당 스택 한도를 1024 B에서 ~14.4 KB로 늘리는 드라이버 syscall로 컴파일합니다. 이로 인해 70-SM RTX 5070 Ti에서는 약 1.34 GiB, 170-SM RTX 5090에서는 약 3.3 GiB가 영구적으로 예약됩니다. 이 예약은 KV 풀과 expert 캐시의 크기가 여유 VRAM을 기준으로 결정된 후에 발생하므로 서버에서 OOM이 발생합니다. 해결 방법: torch 2.14 + Triton 3.8 이상으로 업그레이드하세요. Triton 3.8은 sm_120에서 네이티브 cp.async.bulk.tensor.2d.shared::cta 명령어를 생성하며, 이는 syscall이 필요 없고 스택 예약도 발생시키지 않습니다. 엔진은 이미 이러한 버전을 요구하고 있습니다. 이 문제가 발생하면 오래된 venv 또는 구버전으로 고정된 종속성을 사용하는 소스 빌드가 있는지 확인하고, 업데이트된 requirements로 재설치하세요.

Question 3

chunked prefill 중 발생하는 'expandable_segments: memory mapping failed with OOM' CUDACachingAllocator 경고를 해결하는 방법은?

Accepted Answer

--max-prefill-length를 낮추세요(예: 기본값 8192로 되돌리기). 이 플래그는 서버가 처리할 수 있는 최대 프롬프트 길이가 아니라 chunked-prefill 청크 크기를 설정합니다 — 30k 토큰 요청도 8192에서 문제없이 처리되며, 단지 더 많은 청크로 분할될 뿐입니다. 경고가 발생하는 이유는 큰 청크가 DSA/MLA 어텐션 경로에서 큰 일시적 activation 피크를 만들기 때문입니다(12288 토큰 청크 기준 약 3.8 GiB에 해당하는 여러 개의 활성 [T, 64, 512] bf16 텐서). 이로 인해 torch의 expandable-segments 할당자가 메모리 매핑을 재시도하게 됩니다. 각 재시도 버스트는 디바이스 전체 동기화 스톨(1.5~5초)을 유발하지만 자동으로 해결되므로, 요청은 정확성에 영향 없이 계속 완료됩니다. 참고: overlap 모드에서 쓰기에 MoE expert 캐시가 너무 작은 경우, 청크당 full-expert PCIe 복사 비용이 더 많은 토큰에 걸쳐 분산되므로 더 큰 청크가 전체적으로 더 빠를 수 있습니다 — 따라서 스톨 없는 로그와 최대 prefill 처리량 사이에는 트레이드오프가 있습니다.

Question 4

torch 2.14.1로 업그레이드할 때 발생하는 sglang-kernel 로드 오류와 'C++20 or later compatible compiler is required' 오류를 해결하는 방법

Accepted Answer

sglang-kernel은 libtorch C++ 심볼을 링크하고 빌드된 torch 버전에 대해서만 로드되므로 버전을 함께 고정해야 합니다: `sglang-kernel==0.4.9`는 `torch>=2.14,<2.15`와 함께 사용하세요(0.4.5는 torch 2.11 빌드입니다). torch 2.14에는 Triton 3.8이 필요하므로 `triton>=3.6,<3.9`와 `torchvision>=0.29,<0.30`을 설정하세요. torch 2.14.1의 PyPI wheel은 cu130 빌드이며 `cuda-toolkit==13.0.3`을 요구하므로 installer/uv 인덱스를 cu130으로 유지하세요. C++ 확장을 컴파일하는 경우 setup.py에서 `-std=c++17`을 `-std=c++20`으로 변경하세요. 그렇지 않으면 torch 2.14 헤더가 `#error C++20 or later compatible compiler is required` 오류와 함께 실패합니다. 참고: Triton 3.8은 일부 커널의 성능을 저하시킵니다(예: NVFP4 MoE decode용 `_decode_nvfp4_marlin_kernel`이 40-130% 느려짐). 8개의 성능 저하는 후속 PR에서 수정되었으며, 나머지는 end-to-end 기준 ~4% 미만의 영향을 미칩니다.

Question 5

FreeToken으로 서빙할 때 Qwen3.5/Qwen3.6 RMSNorm 가중치의 bf16 정밀도 손실을 해결하는 방법은?

Accepted Answer

버그: FreeToken의 `models/qwen3_5_moe` 로더는 로드 시점에 Gemma 스타일 `+1`을 bf16 체크포인트 텐서에 합산하여(`tensor + 1.0`) `bf16(1 + w)`를 저장합니다. 반면 Transformers는 제로 중심 가중치를 유지하고 `(1 + w)`를 fp32로 계산하므로, 이 합산 과정에서 작은 가중치가 반올림으로 소실됩니다. `nvidia/Qwen3.6-35B-A3B-NVFP4`에서는 171,008개의 norm 승수 중 77.4%가 변경되며, 평균 상대 오차는 1.4e-3, 최대는 3.9e-3(bf16 ulp의 절반)으로, 모든 토큰에서 `input_layernorm`, `post_attention_layernorm`, `self_attn.q_norm`, `self_attn.k_norm` 및 최종 `model.norm`에 영향을 미칩니다. 수정 방법: 원래의 제로 중심 bf16 가중치 `w`를 그대로 저장하고(로더에서 1을 더하지 마세요) 런타임에 fp32로 `+1`을 적용하세요. 기존 `GemmaPlusOneRMSNorm` 레이어(`layers/norm.py`)를 사용하거나 `sgl_kernel.rmsnorm(x, bf16(1+w))` 대신 `sgl_kernel.gemma_rmsnorm(x, w)`를 호출하면 됩니다. 이는 `qwen4_exp` 로더에서 이미 사용 중인 방식과 일치하며, fp64 레퍼런스와의 bf16 일치율을 99.999%로 복원합니다(잔여 오차는 단순한 출력 반올림이며, rel L2는 이전 3.0e-3에서 ~1.7e-3으로 감소).

FreeToken

주요 기능

할 수 없는 것

사용 사례

태그

시작하기

유사한 프로젝트