qwenGPU 공개, Qwen3.5-2B 가중치 안에서 CUDA 실행하는 실험 모델

개발자 Vixhal이 Qwen3.5-2B의 가중치 안에 GPU를 구현한 실험 모델 qwenGPU를 MIT 라이선스로 GitHub에 공개했다. 이 모델은 자신의 순전파 과정에서 실제 CUDA 코드를 실행한다. 사전 빌드된 체크포인트는 제공되지 않아서, 써 보려면 저장소의 안내대로 직접 만들어야 한다.^1,2

실행에는 Python 3.11 이상, uv, just와 여유 메모리 약 16GB가 필요하다. Apple 실리콘, NVIDIA GPU, CPU에서 돌릴 수 있고 양자화 빌드는 없다. 고급 모드까지 쓰려면 NVPTX 백엔드를 갖춘 clang이 더 필요하다. just build 명령이 Qwen3.5-2B를 내려받아 GPU를 컴파일해 넣는다.^1,2

GPU를 가중치에 넣은 방식

qwenGPU는 Qwen3.5-2B 원래 24개 층 앞에 컴파일된 45개 층을 놓아 모두 69개 층이 된다. 45개 층은 학습하지 않았고, 컴파일러가 가중치를 전부 계산했다. 기반 모델의 컨텍스트 창은 262,144토큰이다.^1,3

이 GPU는 4레인 SIMT 머신으로, 레인마다 레지스터 15개, 4단 파이프라인, 32블록 프로그램 ROM, 126워드 전역 메모리를 갖는다. Qwen3.5의 기존 층 종류인 Gated DeltaNet 선형 어텐션과 SwiGLU MLP만으로 만들었다. 정수 연산, 조건 비교, 갈라지는 분기(divergent branches), 반복문, 공유 메모리, atomicAdd, __syncthreads(), 워프 셔플을 지원한다.

Gated DeltaNet의 순환 상태가 메모리 역할을 한다. 회로의 헤드에서는 게이트를 열어 두고 쓰기 강도를 1로 두며 모든 키를 원-핫으로 만든다. 그러면 갱신마다 상태의 한 행이 정확히 덮어써지고, 원-핫 쿼리로 그 행을 읽어 오므로 각 행이 메모리 셀이 된다. 이 상태에 프로그램 ROM, 레지스터(레인당 워프별 16개, 하나는 항상 0), 전역 메모리가 들어간다.²

CUDA 실행은 호스트와 모델이 나눠 맡는다. 호스트는 실제 GPU의 NVRTC와 드라이버처럼 clang으로 CUDA를 PTX로 컴파일하고, PTX를 기계 블록으로 바꿔 프로그램과 입력 배열을 토큰으로 모델에 넣는다. 이후 실행은 모델이 토큰을 하나씩 쓰며 진행한다. EXEC(b, w) 토큰은 블록 b를 워프 w에 대해 순전파 한 번으로 실행하고, 뒤따르는 VAL 토큰이 레지스터 결과를 싣는다. 메모리 쓰기는 MEM 레코드로 돌아오며, 호스트는 순환 상태에서 결과를 읽을 뿐 계산하지 않는다.

시연 영상에서는 [1, 2, 3]과 [4, 5, 6]의 내적 32를 모델이 토큰 47개로 실행했다. qwenGPU는 Divyansh Agrawal의 Qwenputer가 만든 트랜스포머 컴파일러를 라이브러리로 가져다 쓰고, 그 위에 GPU와 드라이버, 토큰 형식을 새로 만든 것이다.^1,2

개발자가 밝힌 측정 결과

수치는 모두 개발자가 Apple 실리콘의 MPS에서 BF16으로 직접 잰 값이다. 개발자의 실행에서는 모든 실행이 --check 재생 기준으로 정확했다. 이 옵션은 각 실행을 참조 인터프리터와 대조해, 모델이 쓴 토큰이 인터프리터의 토큰과 하나씩 일치하는지 확인한다. 이 인터프리터는 실제 NVIDIA Tesla T4와 테스트 커널 49개 모두에서 결과가 일치했다.^1,2

텍스트 능력은 거의 그대로다. 회로는 텍스트 토큰에서 작동하지 않게 막아 두었고 Qwen의 기존 가중치도 바뀌지 않아서, 합친 모델의 다음 토큰 분포는 원본과 KL 약 7e-4 이내다.²

채팅은 두 모드로 계산한다. 일반 모드에서는 Qwen이 [1, 2, 3] @ [4, 5, 6] 같은 식을 쓰면 kernels/의 기성 CUDA 커널이 각 연산을 실행하며, 표본 채팅 질문 24개 중 20개에 GPU로 계산한 숫자로 답했다. 이는 정답률이 아니라 GPU로 계산한 숫자로 답한 비율이다. 고급 모드에서는 Qwen이 시스템 프롬프트의 CUDA 참조를 보고 커널을 직접 쓰며, 벡터·행렬 질문 5개 중 4개에서 첫 시도에 맞는 커널을 썼다.^1,2

개발자는 2B 모델이 쓰는 CUDA가 아직 믿을 만하지 않다며, 결과를 신뢰해야 하면 일반 모드를 쓰라고 안내했다. 한 번의 실행에는 생성 토큰이 수십에서 수백 개 들고, 속도는 MPS에서 초당 10~15토큰이다.²

MarkTechPost에 따르면 Qwenputer는 호스트가 펌웨어만 올리고 계산은 하지 않으며, Percepta는 2026년 3월에 WebAssembly 인터프리터를 트랜스포머에 넣었다. 이 매체는 속도 수치가 하드웨어와 토큰 종류가 달라 직접 비교할 수 없다고 평가했다.¹

라이선스는 코드가 MIT이다. qgpu.merge로 만든 체크포인트는 설정을 넓히고 GPU 층을 더한 수정 Qwen3.5 모델이라, 가중치에는 Qwen의 Apache-2.0 라이선스가 적용된다.^1,2