RTX 5090 전용 NInfer를 DGX Spark GB10에 포팅해봤다 — 27B는 졌고, 35B MoE는 달랐다

2026년 9월 20일업데이트 2026년 9월 20일약 7분

작성 DevSnack Lab원문·측정 조건·검수 범위는 본문에 기록

NInferDGX SparkGB10CUDALLMMoEbenchmark

NInfer라는 흥미로운 프로젝트를 발견했다.

Qwen3.5 계열 Dense/MoE 모델을 위한 C++/CUDA 추론 엔진인데, 범용성을 넓게 가져가는 llama.cpp나 vLLM과는 방향이 다르다. 지원 GPU와 모델을 좁히는 대신 CUDA Graph, 전용 quantization, MTP/DFlash, paged KV cache 같은 최적화를 적극적으로 사용한다.

문제는 지원 대상이 사실상 RTX 5090 하나라는 점이었다.

빌드 타깃은 sm_120a, 내부에도 5090의 170 SM을 전제로 한 튜닝이 들어가 있었다.

내가 가진 장비는 NVIDIA DGX Spark다.

GB10, compute capability 12.1, 48 SM, ARM64, 128GB unified memory.

그래서 궁금해졌다.

이 엔진을 GB10용으로 포팅하면 실제로 빨라질까?

결과부터 말하면 조금 예상 밖이었다.

27B Dense: 포팅은 성공했지만 별 이득이 없었다.
Qwen3.6-35B-A3B MoE: 상황이 완전히 달랐다.
동일 프롬프트 기준 prefill은 약 3.3배, non-spec decode는 약 1.8배.
MTP3를 사용하면 98~111 tok/s까지 나왔다.

단, 비교한 양자화 포맷은 동일하지 않다. 이 글의 결과는 순수 커널 비교가 아니라 각 런타임에서 실제로 사용할 수 있는 구성의 시스템 비교다.

일단 GB10에서 빌드가 될까

테스트 환경은 다음과 같다.

항목	환경
Device	NVIDIA DGX Spark
GPU	GB10
Compute Capability	12.1
SM	48
CPU Architecture	aarch64
Unified Memory	128GB LPDDR5X
Memory Bandwidth	273 GB/s
Driver	580.178.04
CUDA	13.0 / nvcc 13.0.88

처음 확인한 것은 sm_121a 자체가 현재 CUDA 환경에서 빌드되는지였다.

간단한 CUDA kernel을:

nvcc -arch=sm_121a

로 컴파일하니 정상적으로 통과했다.

그다음 NInfer에서 사용하는 NVFP4 관련 mma.sync를 따로 떼어 테스트했다.

여기서는 잠깐 잘못된 결론을 내렸다.

-arch=sm_121a로 컴파일하면:

Instruction 'mma with block scale' not supported on .target 'sm_121'

라는 에러가 나왔다.

처음에는 "GB10에서 이 명령을 지원하지 않는 건가?"라고 생각했다.

그런데 sm_120a에서도 똑같이 실패했다.

원인은 GPU가 아니라 컴파일 옵션이었다.

NInfer가 실제로 사용하는 것처럼 virtual architecture에도 a suffix가 붙도록:

--generate-code=arch=compute_121a,code=[compute_121a,sm_121a]

형태로 컴파일하자 정상적으로 통과했다.

즉 NVFP4 명령 자체가 GB10에서 막히는 것은 아니었다.

첫 번째 포트

가장 먼저 수정한 것은 CMake였다.

기존 NInfer는 sm_120a 이외의 CUDA architecture를 거부한다.

이를:

120a
121a

두 architecture를 허용하도록 수정했다.

GB10에서는 다음처럼 빌드한다.

cmake -S . -B build -G Ninja \
  -DCMAKE_BUILD_TYPE=Release \
  -DCMAKE_CUDA_ARCHITECTURES=121a

cmake --build build -j

총 437개 target이 정상적으로 빌드됐고, cuobjdump로 생성된 GPU object가 sm_121a인지도 확인했다.

하지만 첫 추론은 실패했다.

Qwen3.5 family runtime requires compute capability 12.0

빌드뿐 아니라 runtime startup 코드에도 12.0만 허용하는 체크가 하나 더 있었다.

이 부분을 12.0과 12.1 모두 허용하도록 수정했다.

이제 모델이 올라갔다.

첫 실험: Qwen3.8-27B NVFP4

처음 테스트한 모델은 Qwen3.8-27B NVFP4였다.

19.0GiB 정도의 weight가 정상적으로 로드됐고 일반 decode도 동작했다.

MTP3도 정상 작동했다.

--spec mtp
--draft-tokens 3
--lm-head-draft

결과는 약:

21.5 tok/s

OpenAI 호환 서버로 실행했을 때는 짧은 요청에서 약 28 tok/s까지 나왔다.

CUDA Graph도 정상적으로 준비됐다.

GPU utilization도 추론 중 95% 수준이었고 clock은 약 2.5GHz까지 올라갔다.

처음에는 생각했다.

"포팅은 됐는데 뭔가 최적화가 빠진 것 아닐까?"

실제로는 그냥 메모리 대역폭 문제였다

Nsys로 짧은 추론을 프로파일링해보니 실행 시간의 약 90%가 GEMM 계열이었다.

GDN recurrent나 gating, attention은 상대적으로 작았다.

그리고 단순하게 계산해보면 이유가 보인다.

NInfer NVFP4 artifact는 약 19.7GB다.

GB10의 메모리 대역폭은 273GB/s다.

1 token을 생성할 때 전체 weight traffic이 지배적이라고 단순화하면:

19.7 GB / 273 GB/s ≈ 72 ms

정도가 된다.

약 13.9 tok/s 수준의 1차 bandwidth roof estimate다.

실제 non-spec decode는 약 10.6 tok/s였다.

즉 커널 몇 개를 조금 더 빠르게 만든다고 갑자기 2~3배 뛰기 어려운 영역이었다.

그리고 이미 GB10에서 사용하고 있던 llama.cpp 쪽에는 더 작은 quantization이 있었다.

Ridge 3.7bpw GGUF는 약 12.6GB였고 MTP를 사용하면 약:

30 tok/s

가 나왔다.

반면 NInfer NVFP4는:

21~28 tok/s

였다.

27B Dense에서는 NInfer를 GB10에 포팅할 실질적인 이유가 없었다.

포팅 자체는 성공했지만 성능상으로는 패배였다.

그런데 35B MoE에서는 결과가 달라졌다

Dense 모델과 MoE 모델의 전체 경로 실행과 선택적 expert 라우팅을 비교한 개념 일러스트 — Dense는 전체 경로를, MoE는 일부 expert를 선택적으로 활성화한다는 차이를 추상화한 이미지입니다.

다음으로 Qwen3.6-35B-A3B를 테스트했다.

35B라는 이름만 보면 27B보다 더 느릴 것 같지만 이 모델은 MoE다.

전체 parameter는 35B지만 inference에서 활성화되는 parameter는 훨씬 작다.

NInfer 공식 groupwise-int artifact를 사용했다.

그리고 MTP3를 켰다.

결과:

98~111 tok/s

MTP acceptance는 프롬프트에 따라 약 46~56%였다.

일반 non-spec decode도:

약 65~69 tok/s

가 나왔다.

이 시점에서 NInfer GB10 포트의 의미가 생겼다.

처음에는 prefill이 느리다고 생각했다

그런데 처음 측정했을 때 prefill이 이상했다.

약 300 tok/s 정도가 찍혔다.

기존 llama.cpp 결과는 1700 tok/s 수준이었다.

"decode는 빠른데 prefill은 처참하네."

라고 생각했다.

그런데 이 비교는 틀렸다.

NInfer 쪽 300 tok/s 결과는 겨우 65 token 정도의 짧은 prompt였다.

조건 자체가 달랐다.

동일한 긴 프롬프트로 다시 측정했다.

약 6,800 token짜리 prompt를 사용했다.

결과:

NInfer default chunk 1024
→ 약 4,350 tok/s

이미 llama.cpp보다 훨씬 빨랐다.

GB10에서는 prefill chunk 4096가 더 빨랐다

NInfer upstream 기본 prefill chunk는 1024다.

이 값은 RTX 5090 환경에서 정해진 기본값이다.

GB10에서도 최적이라는 보장은 없다.

그래서 chunk size를 바꿔가며 측정했다.

Prefill chunk	Throughput
1024	4.35k tok/s
2048	4.89k tok/s
4096	5.31k tok/s
full prompt (~6912)	5.21k tok/s

1024에서 4096으로 바꾸는 것만으로 약 22%가 올라갔다.

4096에서 필요한 workspace도 약 417MiB였다.

128GB unified memory를 가진 GB10에서는 크게 부담되는 크기가 아니다.

그래서 현재 GB10 포크에서는 기본값 자체를 바꾸지는 않고 다음 옵션을 권장한다.

--prefill-chunk 4096

5090용 upstream 동작에는 영향을 주지 않기 위해서다.

같은 프롬프트로 llama.cpp와 비교

최종적으로 약 6,810 token짜리 동일한 프롬프트를 사용해 A/B 테스트했다.

비교 대상은:

llama.cpp
Qwen3.6-35B-A3B Q8_0

그리고:

NInfer GB10 port
Qwen3.6-35B-A3B groupwise-int

였다.

결과는 다음과 같다.

	llama.cpp Q8_0	NInfer groupwise-int
Prefill	1,612 tok/s	5,310 tok/s
Decode, non-spec	36.8 tok/s	~65 tok/s
Decode, MTP3	-	98~111 tok/s

관측된 시스템 성능 차이는 대략:

Prefill       3.29x
Decode        1.77x non-spec
Decode        최대 약 3x with MTP3

였다.

하지만 이 숫자를 "NInfer 엔진이 llama.cpp보다 3배 빠르다"라고 해석하면 안 된다.

양자화 포맷이 다르다.

llama.cpp는 Q8_0이고 NInfer는 자체 groupwise-int artifact다.

따라서 이 테스트는 커널 하나나 inference engine 자체만 격리해 비교한 benchmark가 아니라:

GB10에서 현재 실제로 사용할 수 있는 두 inference configuration의 시스템 비교

다.

그래도 실사용 입장에서는 꽤 의미 있는 차이다.

왜 Dense에서는 졌는데 MoE에서는 이겼을까

GB10의 특성을 생각하면 어느 정도 설명이 된다.

GB10은 128GB unified memory라는 매우 큰 용량을 제공하지만 메모리 대역폭은 273GB/s다.

Dense 모델은 token 하나를 생성할 때 사실상 전체 모델 weight를 계속 읽어야 한다.

그래서 27B에서는 작은 quantization을 사용한 llama.cpp가 유리했다.

19.7GB를 읽는 엔진보다 12.6GB를 읽는 엔진이 유리한 것은 자연스럽다.

반면 MoE는 다르다.

Qwen3.6-35B-A3B는 모든 expert를 매 token마다 계산하지 않는다.

활성 parameter가 작기 때문에 Dense 27B와 같은 방식으로 전체 35B weight가 매 decode step의 직접적인 병목이 되지 않는다.

여기에 NInfer의:

전용 sparse MoE kernel
CUDA Graph
MTP speculative decoding
모델별 native execution path

같은 최적화가 겹치면서 GB10에서도 효과가 나타났다.

결국 GB10에서 이 포트의 핵심은:

Dense 모델을 빠르게 만드는 것이 아니라 MoE inference를 빠르게 만드는 것

에 가까웠다.

170 SM 하드코딩도 GB10에 맞게 수정했다

GPU 실행 자원과 launch tuning을 GB10에 맞게 조정하는 과정을 표현한 개념 일러스트 — 개념 이미지: 실제 GB10의 SM 배치나 하드웨어 토폴로지가 아니라 플랫폼 포팅과 launch tuning을 추상화한 삽화입니다.

NInfer는 5090에 매우 특화된 프로젝트다.

소스를 보면 5090의 170 SM을 기준으로 한 tuning 값들이 들어 있다.

GB10은 48 SM이다.

그래서 device SM count를 runtime에 읽는 작은 helper를 추가했다.

cudaGetDeviceProperties()

를 한 번 호출해 SM 수를 캐시하고 다음 경로에 반영했다.

GDN chunked output distribution
small-T attention split cap
sparse MoE prefill grid cap
RoPE wave threshold

단순히 170을 48로 바꾼 것이 아니라, 170 SM에서는 upstream과 동일한 결과가 나오도록 기존 식을 유지했다.

이 변경 전후도 측정했다.

결과는:

Decode MTP3
21.5 → 21.9 tok/s

Prefill
2.54k → 2.63k tok/s

Non-spec
10.6 → 10.5 tok/s

정도였다.

대부분 오차 범위다.

즉 이 수정이 단일 요청 성능을 극적으로 올려준 것은 아니다.

대신 batch>1과 MoE prefill에서 launch sizing을 실제 GPU topology에 맞게 만드는 성격이 더 크다.

서버도 동작했다

CLI만 실행되는 것으로 끝내지 않고 ninfer-serve도 테스트했다.

OpenAI Chat Completions endpoint가 정상적으로 동작했고 thinking output도 보존됐다.

MTP 역시 정상적으로 accept됐다.

동시 요청 2개도 문제없이 처리됐다.

즉 현재 포트는 단순 benchmark binary가 아니라 실제 OpenAI-compatible inference server로 사용할 수 있다.

그래서 포크를 공개했다

포크:

https://github.com/gdevsnack-ai-labs/ninfer-gb10

현재 포크는 upstream NInfer에 다음 GB10 지원을 추가한다.

sm_121a build target
compute capability 12.1 runtime 허용
aarch64 / DGX Spark 검증
runtime SM count 기반 launch sizing
GB10용 빌드 및 실행 문서
실제 27B Dense / 35B MoE benchmark 결과
GB10에서 --prefill-chunk 4096 권장

CUDA 13.1로 업그레이드할 필요도 없었다.

테스트한 DGX Spark의 기본 CUDA 13.0 환경에서 전체 포팅과 추론이 정상적으로 동작했다.

이 포크가 의미하는 것과 의미하지 않는 것

이 결과를 "NInfer가 DGX Spark에서 항상 llama.cpp보다 빠르다"고 정리하면 틀린다.

실제로 27B Dense에서는 그렇지 않았다.

더 가벼운 GGUF를 사용하는 llama.cpp 쪽이 더 빨랐다.

반대로 Qwen3.6-35B-A3B에서는 NInfer 포트가 상당히 좋은 결과를 보였다.

현재 내 결론은 단순하다.

27B Dense
→ llama.cpp + 가벼운 GGUF

35B-A3B MoE
→ NInfer GB10 port

GB10은 5090과 전혀 다른 하드웨어다.

5090은 높은 메모리 대역폭을 갖고 있고, GB10은 상대적으로 낮은 대역폭 대신 128GB의 unified memory를 제공한다.

따라서 5090 전용 엔진을 GB10에서 단순히 "돌아가게" 만드는 것만으로 충분하지 않았다.

어떤 모델 구조가 GB10의 특성과 맞는지를 같이 봐야 했다.

이번 실험에서는 그 경계가 꽤 명확하게 보였다.

Dense에서는 포팅 자체가 성능 이점을 만들지 못했다.

MoE에서는 달랐다.

그리고 처음 예상했던 것보다 차이가 컸다.

현재 상태

현재 포크는 experimental이다.

RTX 5090 upstream 지원을 대체하려는 프로젝트가 아니라, 기존 NInfer 위에 DGX Spark GB10 (sm_121a) 지원을 추가하는 형태다.

앞으로 확인해볼 만한 것은 upstream에 올라와 있는 sparse MoE, prefix cache, fused kernel, 35B NVFP4 관련 최적화들이다.

하지만 우선은 지금 상태를 baseline으로 남겨두기로 했다.

최적화를 더 섞기 전에:

GB10 functional port
+
SM-aware launch tuning
+
35B-A3B measured baseline

을 먼저 고정해두는 편이 이후 결과를 비교하기 쉽기 때문이다.

이번 포팅에서 가장 재미있었던 건 결국 이것이었다.

처음에는 "RTX 5090 전용 CUDA 엔진을 GB10에서도 돌릴 수 있을까?"가 질문이었다.

실제로 돌리고 나니 질문이 바뀌었다.

"GB10에서는 어떤 모델에서 이 엔진의 구조가 의미가 있을까?"

27B Dense에서는 답이 아니었다.

35B MoE에서는 꽤 분명한 답이 나왔다.