qwenGPU 는 SIMT GPU를 가중치에 직접 컴파일하는 실험적 모델로, Qwen3.5-2B에 탑재되었습니다. 개발자 Vixhal(@TheVixhal)이 MIT 라이선스로 GitHub 에 공개했습니다. 이 모델은 자체 순전파 안에서 NVIDIA PTX로 컴파일된 실제 CUDA를 실행합니다. 개발자들에게 이 모델은 하나의 아이디어를 시험합니다: 사전학습된 LLM 내부에서의 정확한 연산, 곁이 아니라 내부에서.
TL;DR
- 크기: Qwen3.5-2B (24 레이어)에 컴파일된 GPU 레이어 45개를 더해 총 69개. 컨텍스트: 262,144 토큰 (베이스 모델).
- 실행 환경: Apple 실리콘, NVIDIA GPU 또는 CPU에서 약 16 GB의 여유 메모리로 실행. 양자화 빌드 없음.
- 성능: 개발자의 실행에서 모든 실행은 정확했음; 텍스트 드리프트는 KL 약 7e-4.
- 최고: 레퍼런스 인터프리터가 실제 NVIDIA Tesla T4와 49개 커널 중 49개에서 일치.
- 최악: Apple MPS에서 초당 약 10~15 토큰; Qwen이 작성한 CUDA는 5개 중 4개 통과.
qwenGPU란 무엇인가?
qwenGPU는 Qwen3.5-2B에 컴파일된 트랜스포머 레이어 45개 Qwen 자체의 24 앞에 배치된 모델입니다. 그 레이어들은 하나도 훈련되지 않았습니다. 컴파일러가 모든 가중치를 계산합니다.
GPU가 어떻게 모델 가중치 안에 들어가는가?
이 GPU는 4-레인 SIMT 기계입니다. 레인당 15개의 레지스터, 4단계 파이프라인, 32-블록 ROM, 126개의 메모리 워드를 갖습니다. 오직 Qwen3.5의 자체 레이어 유형만 사용합니다: Gated DeltaNet 선형 어텐션과 SwiGLU MLP.
핵심 비결은 Gated DeltaNet 순환 상태입니다. 회로의 헤드에서 게이트는 열린 상태로 유지되고 쓰기 강도는 1입니다. 모든 키는 원-핫입니다. 각 업데이트는 정확히 1개의 행을 덮어쓰고, 원-핫 쿼리가 그것을 다시 읽어냅니다. 그래서 각 행이 메모리 셀로 작동합니다.
그 상태가 프로그램 ROM, 레지스터 및 글로벌 메모리를 보관합니다. README는 레인당 16개의 레지스터 행을 나열하며, 그중 1개는 항상 0을 읽습니다. 이는 X 포스트에서 언급된 15개의 사용 가능한 레지스터와 일치합니다.
이 GPU는 정수 연산, 술어, 분기 분기(divergent branches) 및 루프를 지원합니다. 또한 공유 메모리와 atomicAdd, __syncthreads() 워프 셔플도 지원합니다.
CUDA 런치가 진행되는 동안 무슨 일이 일어날까요?
호스트는 NVRTC와 CUDA 드라이버의 역할을 맡습니다. clang으로 CUDA를 PTX로 컴파일한 뒤, PTX를 머신 블록으로 낮춥니다. 프로그램과 입력 배열을 토큰으로 모델에 공급합니다.
그 후로는 모델이 런치 자체를 작성합니다. 각 EXEC(b, w) 토큰은 1회 순전파에서 블록 b 을 실행하며 w 1회의 순전파(forward pass) 내에서 가능합니다. The VAL 단위로 동작합니다. 그 뒤의 MEM 토큰들이 레지스터 결과를 전달합니다. 메모리 쓰기는
레코드로 반환됩니다. 호스트는 순환 상태에서 결과를 읽으며 그 어느 것도 직접 계산하지 않습니다. 데모 영상에서는 [1, 2, 3]과 [4, 5, 6]의 내적으로 32가 반환됩니다. 모델은 그 런치를 47개의 토큰으로 작성했습니다.
qwenGPU 만나기: 가중치 안에 SIMT GPU를 품은 Qwen3.5-2B로, Qwen 자체 레이어 앞에 45개의 컴파일된 레이어가 있으며, Qwen의 Gated DeltaNet 레이어를 레지스터와 메모리로 사용합니다.
— vixhaℓ (@TheVixhal) 2026년 11월 1일
순전파에서 실제 CUDA를 실행합니다. 호스트가 커널을 PTX로 컴파일해 로드하면, 모델이 모든 …을 작성합니다. pic.twitter.com/iM5IETEbIv
GPU가 Qwen의 텍스트 능력을 깨뜨릴까요?
개발자의 측정 기준으로는 그렇지 않습니다. 회로는 텍스트 토큰에서 음소거되고, Qwen 자체 가중치는 변경되지 않습니다. 병합된 모델의 다음 토큰 분포는 원본과 KL 약 7e-4 이내에 머뭅니다.
qwenGPU는 얼마나 정확할까요?
모든 수치는 자체 보고된 것으로, bf16의 MPS를 사용하는 Apple 실리콘에서 측정되었습니다:
- 모든 런치는
--check리플레이 하에서 정확했습니다. - 레퍼런스 인터프리터가 49개 테스트 커널 중 49개에서 NVIDIA Tesla T4와 일치했습니다.
- 일반 모드는 샘플링된 24개 채팅 질문 중 20개에 GPU로 계산된 숫자로 답했습니다.
- 고급 모드는 5개 질문 중 4개에서 첫 시도에 올바른 커널을 작성했습니다.
일반 모드에서는 Qwen이 준비된 커널을 위한 다음과 같은 표현식을 작성합니다: [1, 2, 3] @ [4, 5, 6] 고급 모드에서는 Qwen이 CUDA 커널 자체를 작성합니다. README는 2B 모델의 CUDA는 아직 신뢰할 수 없다고 경고합니다.
qwenGPU는 Qwenputer 및 Percepta와 어떻게 비교됩니까?
qwenGPU는 다음의 transformer 컴파일러를 사용합니다: Qwenputer 이는 Divyansh Agrawal이 만든 것으로, Qwenputer는 32비트 RISC-V CPU를 Qwen3.5-4B로 컴파일합니다. Percepta 는 2026년 3월에 2D attention head를 갖춘 transformer에 WebAssembly 인터프리터를 구축했습니다.
| 특징 | qwenGPU | Qwenputer (RISC-V) | Percepta |
|---|---|---|---|
| 기반 모델 | Qwen3.5-2B, 24 layers | Qwen3.5-4B | 2D attention head를 갖춘 커스텀 transformer |
| 컴파일된 머신 | 4-lane SIMT GPU | 32비트 RISC-V (RV32IM + Zba + bignum) | WebAssembly 인터프리터 |
| 실행 가능한 프로그램 | CUDA, PTX에서 lowered됨 | C 및 MicroPython | C, WebAssembly로 컴파일됨 |
| 컴파일된 레이어 | 45 | 11 | 공개되지 않음 |
| 상태가 저장되는 위치 | Gated DeltaNet state | Gated DeltaNet state | 컨텍스트 내 실행 trace |
| 호스트의 역할 | CUDA를 PTX로 컴파일하고, 토큰을 로드하며, 결과를 읽음 | 펌웨어를 로드함; 연산은 수행하지 않음 | 공개되지 않음 |
| 정확성 검사 | 인터프리터가 49개 커널 전체 중 49개에서 Tesla T4와 일치했습니다 | 재생된 약 230,000개 포지션에서 불일치 0건 | 수백만 회의 신뢰할 수 있는 스텝을 주장 |
| 신고된 속도 | 10~15 tokens/s (Apple MPS) | 약 3,200 trace tokens/s (L4, Triton) | 초당 30,000+ 토큰 (CPU) |
| 기준 대비 텍스트 드리프트 | KL 약 7e-4 | KL 6e-4 | 해당 없음 |
| 라이선스 | MIT 코드; Apache-2.0 체크포인트 | MIT 코드; Apache-2.0 가중치 | 공개되지 않음 |
속도 수치는 서로 다른 하드웨어와 토큰 유형을 사용하므로 직접 비교할 수 없습니다.
이를 실행하려면 무엇이 필요한가요?
그 빠른 시작 Python 3.11+, uv, just 및 약 16 GB의 여유 메모리가 필요합니다. 고급 모드는 NVPTX 백엔드가 포함된 clang도 필요로 합니다. just build Qwen3.5-2B를 다운로드하여 GPU를 그 안에 컴파일합니다. 사전 빌드된 체크포인트는 공개되지 않습니다. 빌드된 체크포인트는 Qwen의 Apache-2.0 라이선스를 따릅니다.
핵심 요약
- qwenGPU는 4레인 SIMT GPU를 Qwen3.5-2B에 45개의 미학습 레이어로 컴파일합니다.
- Gated DeltaNet 상태는 레지스터, 프로그램 ROM 및 전역 메모리 역할을 합니다.
- 모델이 모든 CUDA 실행 단계를 작성하며, 호스트는 컴파일과 로드만 수행한다.
- 인터프리터가 49개 커널 중 49개에서 Tesla T4와 일치했습니다.
- 초당 10~15 토큰의 속도로, 이것은 런타임이 아니라 연구 데모입니다.
다음을 확인해 보세요 GitHub 저장소, 이 qgpu CUDA 참조, 그 Qwen3.5-2B 모델 카드 그리고 그 원본 X 게시물. 또한 저희를 팔로우해 주세요 Twitter 그리고 저희의 150k+ML SubReddit 에 가입하고 저희의 our Newsletter을 구독하는 것도 잊지 마세요. 잠깐! 텔레그램을 사용하시나요? 이제 텔레그램에서도 저희를 만나볼 수 있습니다.
GitHub Repo, Hugging Face Page, 제품 출시, 웨비나 등을 홍보하기 위해 저희와 파트너십이 필요하신가요? 저희에게 연락하세요
다음 포스트 Meet qwenGPU: Qwen3.5-2B With 45 Compiled GPU Layers That Run CUDA 는 다음에서 처음 게시되었습니다 MarkTechPost.