PocketLLM: 단일 가속기에서 대규모 언어 모델 실행

English | 中文

하나의 가속기—단일 GPU, 엣지 보드, 주머니 속 휴대폰—에서 대규모 언어 모델을 실행합니다. 하나의 프로세스가 하나의 장치를 소유합니다. 체크포인트가 맞지 않으면 더 정량화(quantized)됩니다.

상태: C 엔진은 Qwen3-0.6B의 f16 또는 q4_k_m으로 실행합니다.

; Python 패키지는 모델을 실행하지 않습니다. src/

(libpocketllm.so)는 GGUF를 읽고, 체크포인트 자체의 BPE로 토큰화하며, Qwen3 그래프를 따라가고, f32/f16 가중치와 커널 내부에서 디코딩되고 절대로 확장되지 않는 패킹된 q4_k/q6_k으로 탐욕적으로(greedily) 디코드합니다. 이 과정은 cpu 및 cuda 카드 모두에서 llama.cpp와 토큰별로 검증되었습니다. 1.4 GB의 f16 체크포인트와 그로부터 정량화된 456 MB의 q4_k_m에 대해 각각, 각 백엔드는 구현하는 어텐션 규칙에 따라 테스트되었습니다. 왜냐하면 llama.cpp의 기본 flash-attention 모드와 전체 softmax 모드가 서로 다른 산술 방식을 사용하고 거의 동점인 경우 다른 토큰을 선택하기 때문입니다. Python 패키지는 호스트 측—사양으로서의 커널 ABI, numpy oracle, GGUF 로더, 정량화 디코더, 실행 계층 및 OpenAI-호환 HTTP 인터페이스—이며, 어떤 Python 백엔드도 커널을 구현하지 않습니다: reference를 제외한 모든 백엔드는 세션에서 BackendNotImplementedError가 발생하는 선언입니다. 이들은 정확히 한 지점에서 만납니다:

python/pocketllm/native.py,

두 가지 모두 구동하는 ctypes 브리지입니다—pocketllm run과 pocketllm serve 둘 다 사용합니다—. `pocketllm run --model ckpt.gguf --prompt

가중치 폭(width ladder)은 Q4 → Q2 → IQ2 → IQ1 → ternary 순서로, 모델이 여전히 정확하게 응답하는 가장 낮은 형식에서 멈춥니다. 호스트 오프로드 및 다중 카드 병렬성은 구조상 제외됩니다: ABI에는 랭크(rank), 컬렉티브(collective), 또는 두 번째 장치가 없으며, EngineArgs에는 설정할 tensor_parallel_size가 없습니다.

작동하는 부분, 스텁(stub)으로 존재하는 부분, 그리고 아직 작성되지 않은 부분이 있습니다. 이 표가 솔직한 내용이며, 이 페이지의 나머지 부분은 해당 조각들이 구축될 방향에 대한 설계입니다.

C 코어에서 server/native_backend.py를 통해 구동됩니다. 한 번에 하나의 요청만 처리하며, 배치 처리는 없고 취소 기능도 없습니다.

pocketllm.cli

여섯 가지 명령어(devices, backends, architectures, ops, run, serve) 모두 완료되었습니다.

src/ — C++ 엔진 (libpocketllm.so)

Qwen3-0.6B 모델을 f16 및 GGUF 읽기, BPE 토큰화, 순전파(forward), 그리디 디코드(greedy decode), 그리고 온도/top-k/top-p/min-p 샘플링으로 구동합니다; q4_k_m 및 q4_k/q6_k가 커널에서 디코딩됩니다; 그리디 디코드는 cpu와 cuda 모두에서 llama.cpp와 토큰별로 검사되었습니다 (각 백엔드는 구현하는 어텐션 규칙에 따라 — C 엔진 페이지 참조); 샘플러는 numpy 참조와 토큰별로 검사되었습니다. CPU의 gemm_quant는 스레드 처리되고, AVX2 벡터화되었으며, 활성화를 int8로 양자화하여 llama.cpp 방식과 동일하게 수행합니다 — 이는 쓰레드 풀에서 작업당 하나의 블록을 사용하며, 512×1024 프리필(prefill) 형태에서 양자화기 자체의 속도를 1087 us에서 55.6 us로 단축했습니다; 이 풀은 hardware_concurrency가 아닌 물리적 코어 수(여기서는 44개, 88개가 아님)를 기본값으로 사용하는데, 이는 코어 수를 초과하는 스피닝 워커 풀이 하이퍼스레드 형제들끼리 경쟁하기 때문입니다 — 디코드 시에는 1.3×, 프리필 시에는 1.09×의 성능 향상을 가져옵니다 — 그리고 $POCKETLLM_CPU_PIN=node 또는 =all을 사용하여 자체 친화도(affinity)를 설정할 수 있습니다 (기본값은 비활성화되어 있는데, 풀을 제한하는 것이 이 엔진에 도움을 주는 것보다 llama.cpp에 더 도움이 되기 때문입니다 — C 엔진 페이지 참조); 그 패킹된 점 곱셈(packed dot)은 활성화 행 8개당 한 번씩 가중치 행을 순회하며 (행당 한 번이 아님, $POCKETLLM_CPU_GEMM_RPW=4는 이전 카운트를 선택함), 단일 행 커널과 비트 단위로 동일하며, 행 타일은 4개 행 형태 대비 pp512의 8% 성능 향상을 가져옵니다; '어텐션(attention)' 점 곱셈은 네 개의 레인 폭을 가지며 대체한 스칼라 dot과 비트 단위로 동일합니다. 이것이 짧은 컨텍스트 디코드 속도를 llama.cpp와 동등하게 만들었으며, 그 점 통과는 키 행당 8개의 쿼리 행을 순회합니다 — 이는 인과적 테일(causal tail)을 단일 행 커널에 남긴 하나의 타일 코드 경로로, 결과가 모든 타일 너비에서 dot4와 같도록 쌍으로 매칭되어 있으며, 스칼라 dot과 np.array_equal을 유지합니다. 이 타일의 쌍은...

이제 레지스터는 키마다 재구성되는 대신 타일당 한 번 인터리브(interleaved)됩니다 (비트 단위로 동일하며, 같은 레인과 같은 mul/add 순서를 가지며, 22개의 스레드에서 pp512 대비 1.02–1.08배 및 pp2048 대비 1.08–1.15배의 가치가 있습니다. 이로 인해 점수 계산(score dot)의 키당 vinsertf128은 pp2048 프로파일의 23%를 차지합니다); 로우 타일(row tile)은 어텐션 호출에서 1.27배의 가치를 가지며, 대체한 너비 대비 1.22배입니다. 이로 인해 가중치 합계는 V 행당 여덟 개의 출력 행을 거칩니다 — 이는 호출의 다른 끝단에서도 동일하게 타일링되며, 1부터 512까지 모든 청크에 대한 로우-별 루프와 비트 단위로 정확합니다. 그리고 점수 패스(score pass)는 키당 두 개의 쿼리 헤드(query heads)를 계산합니다 — 이는 KV 헤드를 공유하는 두 헤드이므로, 그룹화된 어텐션 키 행은 두 번 대신 한 번 로드됩니다. 이 페어링은 하나의 __m256의 두 절반으로 표현되어 결과가 비트 단위로 두 개의 dot4 호출을 이루며, 격리된 점수 패스에서는 1.49배, 디코드 형태의 컨텍스트에서 어텐션 호출에서는 2.1배의 가치를 가집니다. 또한 K/V 캐시가 이제 f16이 되었습니다 — 이는 llama.cpp의 -ctk /-ctv 기본값인 너비입니다. 따라서 두 엔진은 동일한 바이트를 스트리밍하고 캐시는 f32 크기의 절반만 차지합니다. 이 캐시는 각 행을 f32로 확장하여 동일한 점수 커널(score kernels)을 실행함으로써 읽히며, 이것이 f16 경로의 전체 정확성 주장입니다: 즉, 확장이 디코드와 프리필 형태 모두에서 비트 단위로 반올림된 f32 캐시를 재현한다는 것입니다. 이는 허용 오차(tolerance)가 아닌 등가성(equality)으로 테스트되었습니다. 확장 및 축소는 256비트로 이루어지며 — 두 개의 128비트 변환 대신 여덟 개당 하나의 vcvtph2ps /vcvtps2ph를 사용합니다 — 이는 비트 단위로 동일하며 (확장은 정확하고, 축소는 같은 방식으로 반올림함), 커널에서 확장 시 1.72–2.57배, 축소 시 1.13–1.15배의 가치를 가지며, kv_row_to_float의 자체 시간이 pp2048에서 5.7%에서 3.9%로 이동하고 약 ~1.02배의 가치를 가집니다.

캐시를 절반으로 줄이는 것은 1024행 컨텍스트에서 디코드 시 tg128로 1.19배, 프리필 시 1.08배의 이득을 가져오며, 256행에서는 캐시가 충분하지 않아 그 가치가 상쇄되지 않습니다. 또한 attention은 두 경로 모두에 플래시 어텐션(flash attention)을 실행합니다. 키(keys)는 점수 행렬을 생성하여 두 번 스캔하는 방식 대신 한 번의 순회 최대값, 분모 및 가중 합계로 처리되는데, 이는 llama.cpp 자체 배포 기본값이 사용하는 방식입니다. 그리고 토큰 하나를 디코드할 때 구간을 log-sum-exp으로 병합된 청크(chunks)로 나누어 8개의 KV 헤드가 유일한 단위가 아니게 합니다. 매칭된 KV 깊이(llama-bench의 tg64는 -d 옵션이 주어지지 않으면 깊이 0에서 실행되는 반면, 이 벤치는 디코딩 전에 프리필을 수행하므로 두 테스트는 -d 512일 때만 비교 가능합니다)로 측정하고 페어링하여 — 양쪽 모두 백투백(back to back)으로 실행되어 동일한 호스트 부하가 모두에 영향을 미치는데, 이는 이 공유 호스트가 지원하는 유일한 통계입니다 — 플래시는 22 스레드에서 디코드 시 1.18배, 44 스레드에서 1.22배의 가치를 가지며, 프리필에서는 0.94배로, 12라운드에 걸쳐 일관성을 보였습니다. 이는 디코드를 llama.cpp 배포 기본값과 동등한 수준으로 만들고, 프리필은 단일 소켓에서 약 1/10 정도 뒤처지게 합니다 (pp512 /tg64 486/63.3 대 22 스레드에서의 547/59.3; 658/67.4 대 918/65.2 at 44). 반면, 두 경우 모두에서 -fa 0보다 우수합니다 (프리필 1.14배, 디코드 1.31배). 이 폴드(fold)의 fp-contract=off는 더 이상 사용되지 않습니다 — 이를 유지했던 상위-2-마진 주장은 현재 part 레이아웃에서 재현되지 않으며, 해당 폴드를 축소하면 (전체 로짓 벡터가 9개 위치에 걸쳐 두 체크포인트 간 바이트 단위로 일치함에도 불구하고) +3.0%의 가치를 지닙니다; C 엔진 페이지를 참조하십시오.

융합된 패스(fused pass)는 토큰 시퀀스를 움직이지 않게 남겨둡니다 (llama.cpp 비교, 증분 대 배치 동등성 및 CPU 대 CUDA 경계 패스는 변경되지 않았음) 그리고 테스트의 바이트 정확도에 비용을 지불합니다 — 가중치 합 타일링(weighted-sum tiling)의 계약은 이제 t22에서 pp512이고 pp2048에서는 +7.6%인 1e-5 스케일 상대 경계로 측정됩니다. 왜냐하면 재스케일링(rescale)을 통해 행의 항들이 한 번에가 아니라 블록으로 만나기 때문이며, 이는 5.96e-07로 측정되었습니다; silu_mul의 쓰레드 그레인(thread grain)은 프리필(prefill)이 아닌 디코드(decode) 모양에 맞춰 크기가 조정되었으며, ctx-512 디코드 토큰의 8%만큼 가치가 있습니다; 프리필의 GEMM 절반은 이미 llama.cpp와 일치합니다 (숫자, 비교 및 왜 거기에 융합 곱셈-덧셈(fused multiply-add)이 있었으면 다른 모델이었는지에 대한 내용은 C 엔진 페이지를 참조하세요); q4_K 가중치도 이제 8열 패널로 읽을 수 있습니다 — llama.cpp 자체의 사전 타일링된 GEMM으로, 행 분할(row-split) 대신 패널 분할(panel-split) 방식으로 포팅되었으며, 매칭되는 패널 GEMV를 통해 디코드가 행 커널에 뒤처지지 않게 합니다 — 이는 쌍으로 측정되고 인터리브되어 1.14×의 가치이며, 기본적으로 pp512이고 단일 소켓 코어에서 tg64의 1.15×입니다 ($POCKETLLM_CPU_REPACK=0은 행 커널을 복원합니다 — 이 트리의 선택자 중 의미가 반전된 유일한 것입니다. 왜냐하면 배포된 기본값은 빠른 것이어야 하기 때문입니다); 여기에서 의도적으로 행 커널과 비트 단위로 동일하지 않은 유일한 커널이 있으며 (그 하위 블록들은 int16으로 누적되는데, 이는 q4_K에 대해서는 합법적이지만 q6_K에서는 오버플로우가 발생하므로, q6_K — 이 체크포인트의 output.weight를 포함하여 — 행 경로에 남아 있습니다), 그리고 전체 스위트는 llama.cpp에 대한 정확한 그리디 토큰 일치(exact greedy-token match)를 포함하여 어쨌든 1001개 통과 / 70개 건너뛰기로 통과합니다; 동일 호스트의 기본 22개 쓰레드에서, 그리고 두 바이너리가 고정된 상태로 같은 세션에서 llama.cpp와 인터리브되어 프리필이 이제 앞서갑니다: pp512 1.037× 및 pp2048 1.034× (각 크기별 6/8 인터리브 라운드 선행, 중앙값 비율), 디코드는 매칭된 KV 깊이에서 두 엔진으로 1.37×입니다 (llama-bench는 -d 512가 필요하며, -d 0에서는 두 가지가 다른 작업을 측정합니다.)

44 스레드에서 pp2048은 1.014배 (7/8 앞섬); pp512는 여전히 뒤처지는 크기(중앙값 0.977배, 2/8)이며, 실행 시간이 가장 짧아 이 공유 호스트의 다른 사용자에게 가장 민감합니다—가장 적게 로드된 라운드는 1.005배였습니다. 이전 파일 기록에서 pp2048이 ~0.99배, 패리티였던 것과 t44는 크게 뒤처지는(0.68배) 수치들은 컨텐디드-윈도우 아티팩트였습니다: 나머지 동일한 스윕 전반에 걸쳐 같은 바이너리 읽기인 pp2048이 1.00–1.13을 기록했습니다. C 엔진 페이지를 참조하세요; build/pocketllm-bench가 이를 측정합니다 |

거기에 main은 없습니다.

시드 커밋 이전의 브랜치 히스토리는 다음과 같습니다: 이 트리는 오펀(orphan) 브랜치에서 재구축되었으며, 이전 것은 legacy로 보존됩니다.

. 코드가 현재 어디에 있는지 참조하세요.

기본 설치는 하나의 의존성입니다. ABI는 stdlib 전용이며, 레퍼런스 백엔드는 numpy이고,
장치 런타임은 선택적 추가 기능입니다—이것이 동일한 휠(wheel)을 CUDA 박스와 휴대폰 모두에 설치할 수 있게 해줍니다.

pip install pocketllm

# 장치 런타임 사용 시
pip install "pocketllm[cuda]" # NVIDIA, torch를 통해
pip install "pocketllm[mps]" # Apple Silicon, torch를 통해

요구 사항: Python >= 3.10. 기본 설치에는 다른 것이 필요하지 않습니다—컴파일러도, CUDA 툴킷도, relic-core도 없습니다.

. 이 트리에는 ext_modules가 없고 빌드 단계도 없습니다: 모든 네이티브 커널은 다른 저장소에 속합니다.

작동하는 사본을 위해:

git clone https://github.com/lvyufeng/PocketLLM.git
cd PocketLLM
pip install -e ".[dev]"

읽기 전용 명령어는 가속기가 없고 torch도 없는 호스트를 포함하여 모든 호스트에서 작동합니다:

pocketllm devices # 이 호스트가 열 수 있는 것, 그리고 나머지가 누락된 것
pocketllm backends # 각 백엔드가 선언하는 것, 사용 가능한지 여부와 관계없이
pocketllm architectures # 이 트리가 구축할 수 있는 모델 구조들
...

devices