GeulPU: 같은 입력의 계산을 재현하는 엔진과 그 한계

배경과 문제

LLM을 로컬 GPU에서 실행할 때 같은 입력과 설정으로도 계산 결과가 달라질 수 있다. 함께 묶이는 요청 수가 바뀌면 커널이 선택하는 타일링이나 부분합을 합치는 순서가 달라지기 때문이다. 이런 차이가 다음 토큰 선택에 영향을 주면 뒤에 이어지는 답도 달라질 수 있다. 생성의 무작위성을 줄이는 설정과 계산 과정의 재현성은 구분할 문제다.

이런 현상은 평가 점수의 재현, 조건 간 비교, 실패 원인을 찾는 데 영향을 준다. 작성자는 같은 GGUF 파일로 llama.cpp의 동시 요청을 1개에서 8개로 바꿨을 때 gsm8k 100문항 중 41개의 생성 결과가 바뀌었다고 보고한다. 이는 답 문장이 달라졌다는 측정이지 정답률이 41% 떨어졌다는 뜻은 아니다.

기존 접근의 한계

기존의 배치 불변 커널은 보통 축약 순서를 고정해서 이 문제를 해결하려 한다. 하지만 부동소수점 덧셈은 결합 법칙이 성립하지 않아, 더하는 순서가 바뀌면 끝자리가 달라진다. 예를 들어 float32 배열을 앞에서부터 차례로 더한 값과 거꾸로 더한 값은 같지 않다.

GeulPU는 순서를 고정하는 대신, 곱의 합을 정수로 정확하게 구하는 방식을 택한다. 정수 덧셈은 결합 법칙이 성립하므로 어떤 순서로 나눠 더해도 결과가 같다. 이 성질을 엔진 전체에 적용하고, 남는 부동소수 연산의 반올림 지점을 모두 문서로 고정한 것이 '정밀도 계약'이다.

연구방법을 단계별로

GeulPU의 정밀도 계약은 활성값, 가중치, 행렬곱, 부동소수 연산, 어텐션, 비선형 함수, 샘플링까지 세부적으로 규정한다. 활성값은 32개 블록마다 2의 거듭제곱 지수 하나와 정수로 표현하고, 정확 모드는 블록 최댓값의 2⁻¹⁴ 격자(int8 두 자리)를, 빠른 모드는 int8 한 자리를 사용한다. 가중치는 GGUF 파일의 정수와 블록 스케일을 그대로 쓴다.

행렬곱에서 프리필은 정수 mma.sync(s8·s8→s32), 디코드는 dp4a를 사용한다. 블록 안의 합은 정확한 정수라 타일링·split·배치와 무관하고, 블록 결과는 오름차순으로 고정된 fma 체인으로 합친다. 모든 f32 연산에는 명시적 .rn 반올림을 적용해 드라이버가 곱셈과 덧셈을 FMA로 합치지 못하게 하고, 근사 명령을 금지한다. 어텐션에서는 KV 캐시도 블록 정수로 저장하고, softmax 분모는 정수 정확합으로 구하며, 64키 타일을 위치로만 정해지는 그룹으로 2단 병합한다.

비선형 함수에서 exp는 Cephes 계수의 fp32 fma 다항식으로, RMSNorm은 256-lane 고정 트리로, RoPE 표는 Python decimal로 정확 반올림한 값을 쓴다. 샘플링과 로그확률은 float64 고정 순서로 계산하고, 난수는 splitmix64(seed, step)를 사용해 배치나 추측 검증 행 수가 난수를 바꾸지 못하게 한다. 이 규칙은 GPU 구현을 NumPy CPU 참조와 대조하기 위한 기준이다. 계약을 정했다는 사실과 실제 구현이 모든 경우에 계약을 지킨다는 검증은 별개다.

결과와 숫자의 의미

개발자의 자체 테스트에 따르면, 같은 GGUF 파일과 RTX 4070 Ti 기준으로 동시 요청 1↔8에서 lm-eval gsm8k 100문항의 문항별 결과가 GeulPU에서는 동일했지만 llama.cpp(b11643)에서는 41문항의 생성이 바뀌었다. 추측 디코딩을 켜고 끄는 경우에도 GeulPU는 토큰·로짓이 동일했지만 llama.cpp(b11721, ngram-simple)는 탐욕 답 8개 중 2개가 변했다. 디스크 프롬프트 캐시 복원과 재계산 비교에서도 GeulPU는 54개 에이전트 요청의 1,728단계가 모두 동일했다.

GPU 교체 실험에서는 정확 모드에서 RTX 3080 Ti·4070 Ti·5070·RX 9060 XT에서 82,673단계의 로짓이 모두 일치했다. 다만 이는 특정 모델(R1-Distill-Qwen-7B)과 수학 32문제에 한정된 결과이며, 모든 GPU·모델·버전을 증명한 것은 아니다. 속도 측면에서는 빠른 모드에서 Qwen2.5-7B Q4_K_M의 디코딩 시간 비가 llama.cpp 대비 0.97~1.06배로 대등했지만, 정확 모드의 프롬프트 시간은 1.16~1.47배로 느렸다.

구체적인 예시

배치 불변성의 문제를 보여주는 간단한 예로, PyTorch에서 2048×4096 행렬 a와 4096×4096 행렬 b를 만들고, a의 첫 행만 계산한 결과(out1)와 2048행을 함께 계산한 뒤 첫 행을 취한 결과(out2)를 비교하면 같은 행인데도 차이가 0이 아닐 수 있다. 이는 GPU 커널이 배치 크기에 따라 다른 축약 순서를 선택하기 때문이다.

GeulPU에서는 활성값을 정수 격자로 옮기면 정수 합은 순서와 무관해진다. 예를 들어 float32 배열을 2¹³을 곱해 정수로 반올림한 뒤 합하면, 원래 순서와 거꾸로 순서의 합이 같아진다. 이 원리를 행렬곱과 어텐션에 적용해 배치가 달라도 계약상 결과를 재현하려는 설계다. 다만 같은 GGUF 파일·계산 모드·컨텍스트 및 RoPE 설정·시드를 맞춰야 비교할 수 있고, 정수 격자로 옮기는 과정 자체의 반올림은 남는다.

한계와 오해하기 쉬운 점

결정성은 재현 도구이지 정답·환각 방지 보장이 아니다. 작성자의 별도 연구 초안은 두 Qwen 계열 7B 모델, Q8_0, 탐욕 디코딩, RTX 4070 Ti에서 한국어 696개와 수학 600개 문제를 비교했다. GeulPU 대비 llama.cpp(b11669)의 다른 문자 섞임은 새로 생김 36·사라짐 35, 수학 오답 또는 미완료는 25·28이었다. 순증가가 통계적으로 뚜렷하지 않았지만, 중간 크기의 악화까지 배제한 동등성 증명은 아니다. 반복 검정도 표본이 부족했고, 판정기는 작성자가 작은 표본으로 점검했다. 이 초안은 동료심사를 마친 논문으로 확인되지 않았다.

같은 연구의 반복 방향 +0.0039nat는 토큰 후보의 로그 승산 변화이며, 실제 반복 답변 발생률이 0.4%포인트 늘었다는 뜻이 아니다. 정수 부분합이 정확해도 활성값 반올림과 비선형 함수 근사 손실은 존재한다. 또한 포럼 글 이후의 v0.10.2는 여유 VRAM에 따른 YaRN 설정, 교체한 모델의 캐시, 14B 모델의 RMSNorm 계수 오류를 고쳤다고 명시한다. 따라서 이전 버전까지 무조건 CPU 참조와 같은 계산이었다고 볼 수 없다.

속도 표는 모델·모드·문맥 길이와 선택 기능에 따라 다르다. v0.10.0은 Qwen3의 F2 품질 기준 미달을 기록하지만 현재 main 문서는 F2B를 다루므로 한 버전의 기능으로 합칠 수 없다. 배포 실행 파일은 Windows x64이며 Linux는 CUDA 중심의 소스 빌드, macOS는 미지원이다. MoE·Gemma·멀티모달·F16/BF16 가중치도 미지원이고, AMD Windows는 Q8_0/Q4_0 정확 모드로 제한된다. GPU 시험은 저자와 지인 장비의 자체 측정이며 이번 검토에서 독립 실행해 재현하지 않았다.

실제로 어떻게 활용할까

GeulPU는 평가 재현, 로그확률·로짓 해시 비교, 짝 비교, 실패 재현과 반사실 분기, GPU 없이 검증 등에 활용할 수 있다. lm-evaluation-harness가 local-completions로 붙고, 작성자는 시험한 조건에서 동시 요청 수와 CUDA/Vulkan을 바꿔도 문항별 결과가 같았다고 보고한다. OpenAI 호환 API에서 단계마다 로짓 행의 64비트 해시를 받아 다른 PC와 비교하면 재현 여부를 간편하게 점검할 수 있다. 다만 짧은 해시가 같다는 것만으로 원래 로짓 배열의 완전한 동일성을 수학적으로 증명하는 것은 아니다.

디버거로 두 양자화를 요청별로 비교하면 같은 답인지, 처음 갈라지는 단계가 어디인지, 같은 답의 로그확률 차이가 얼마인지 알 수 있다. 실행 기록과 재현 파일은 CPU 참조로 확인할 수 있어 리뷰어나 공동 연구자에게 파일 하나만 넘기면 된다. 다만 현재는 연구 단계의 도구이며, 지원 범위와 검증 규모의 한계를 고려해야 한다.