하드웨어 인식 무손실 압축을 활용한 빠르고 메모리 효율적인 LLM 추론

1. 서론 대형 언어 모델(LLM)은 수십억 개의 파라미터를 보유하고 있어 GPU 메모리와 메모리 대역폭이 서비스 단계에서 가장 큰 제약이 된다. 기존의 손실 있는 압축(양자화, 프루닝)은 정확도 저하 위험이 있어, 정확성을 보장해야 하는 실시간 서비스에는 부적합하다. 반면, 손실 없는 압축은 비트‑정확성을 유지하지만, 전통적인 엔트로피 코덱이 생성하는 가변 길이 비트스트림은 GPU의 SIMT 실행 모델과 충돌해 디코딩 단계에서 심각한 성능 저하를 초래한다. 논문은 이러한 구조적 불일치를 해소하고, 압축 이득을 실제 추론 가속으로 전환하기 위한 공동 설계 접근법을 제시한다. 2. 배경 LLM 추론은 주로 두 단계(프리‑필 및 디코드)에서 대규모 행렬 곱(GEMM)을 수행한다. BF16 포맷은 1비트 부호, 8비트 지수, 7비트 맨티사로 구성되며, 현재 대부분의 GPU 텐서 코어가 BF16 연산을 네이티브로 지원한다. 텐서 코어는 warp‑단위로 고정 크기 매트릭스 조각을 동시에 연산하므로, 입력 데이터가 동일한 형식·길이로 제공될 때 최고의 효율을 낸다. 3. 압축 가능성 분석 저자들은 Llama‑3‑8B‑Instruct, Mistral‑24B‑Instruct‑2501, Qwen2.5‑32B‑Instruct 등 최신 모델의 BF16 가중치를 분석하였다. 지수 비트는 8비트 전체 중 2.5‑2.7비트의 엔트로피만을 갖고, 상위 3개의 지수가 전체의 67% 이상, 상위 7개가 95% 이상을 차지한다. 또한 99.6%의 행렬에서 이 상위 7개의 지수가 연속적인 정수 구간을 이룬다. 이는 “지수 연속성”이라는 구조적 특성으로, 고정 길이 인코딩이 가능함을 의미한다. 4. 기존 접근법의 한계 DFloat11(휴먼 코딩)과 DietGPU(ANS)와 같은 기존 무손실 압축은 가변 길이 심볼을 디코딩하기 위해 비트 포인터를 동적으로 이동하고, 각 스레드가 서로 다른 길이의 심볼을 처리하면서 워프 내 분기와 스레드 정체가 발생한다. 실험 결과, L40S GPU에서 이러한 디코더는 메모리 대역폭 활용도가 43.7%~76.5%에 불과했다. 또한, 압축된 가중치를 완전히 복원해 전역 메모리에 저장한 뒤 GEMM에 전달하는 ‘디커플드 파이프라인’은 중복 메모리 트래픽을 유발해 Compute Intensity(CI)를 크게 감소시킨다. 5. ZipServ 설계 5.1 Tensor‑Core‑Aware Triple Bitmap Encoding(TCA‑TBE) - 각 가중치를 (베이스 지수, 오프셋 비트맵, 맨티사) 세 부분으로 고정 길이 16비트 형태로 재구성한다. - 베이스 지수는 각 행렬에 대해 가장 빈번한 연속 지수 구간의 첫 번째 값을 저장하고, 오프셋 비트맵은 7비트(0~6) 범위 내에서 실제 지수를 표현한다. - 맨티사는 그대로 7비트로 유지한다. - 이렇게 하면 모든 스레드가 동일한 연산 흐름(비트 마스크, 시프트, OR)으로 병렬 디코딩이 가능해 SIMT 효율을 회복한다. 5.2 ZipGEMM 커널 - 기존 GEMM 파이프라인을 그대로 사용하되, 입력 가중치를 global memory에서 읽어올 때 압축 블록을 128‑bit 워드 단위로 로드한다. - 워프 내 각 스레드는 로드된 워드에 대해 비트 연산으로 베이스 지수와 오프셋을 결합해 BF16 지수를 복원하고, 맨티사와 결합해 최종 BF16 값을 레지스터에 저장한다. - 복원된 값은 즉시 Tensor Core의 mma.sync.m16n8k16 연산에 투입되며, 중간 버퍼가 존재하지 않는다. - 메모리 접근은 압축 비율에 비례해 감소하고, 연산은 기존 cuBLAS와 동일한 워프‑레벨 스케줄링을 유지한다. 5.3 구현 세부 사항 - 압축 단계는 사전 학습된 모델에 대해 한 번만 수행되며, 각 레이어별 베이스 지수를 메타데이터로 저장한다. - 디코딩 시 필요한 베이스 지수 테이블은 공유 메모리에 캐시되어 워프 내 모든 스레드가 빠르게 접근한다. - 비트 연산은 CUDA의 __brev, __popc 등 하드웨어 가속 명령을 활용해 레이턴시를 최소화한다. 6. 평가 - 모델: Llama‑3‑8B, Mistral‑24B, Qwen2.5‑32B, Gemma‑3‑2B 등 4가지 대표 LLM. - 하드웨어: RTX 4090, L40S, RTX 5090. - 압축률: 25%~30% 감소, 평균 27.8% 압축. - 커널 성능: ZipGEMM이 cuBLAS 대비 1.68×~2.21×, DFloat11 대비 2.9×~5.5× 향상. - 전체 파이프라인: vLLM 대비 평균 1.22×, 최악‑사례 1.07×, 최고‑사례 1.38× 가속. - 메모리 대역폭 사용량은 압축 비율에 비례해 감소했으며, Roofline 분석에서 CI가 압축되지 않은 GEMM보다 50% 이상 향상된 것으로 확인되었다. 7. 논의 및 한계 - 현재 설계는 BF16 지수의 낮은 엔트로피에 특화되어 있어, FP16·FP32 등 다른 포맷에 대한 적용 가능성은 추가 연구가 필요하다. - 다중 GPU 환경에서 압축된 가중치를 네트워크를 통해 전송할 경우, 압축‑전송‑디코딩 파이프라인 전체 최적화가 요구된다. - 압축 메타데이터(베이스 지수 테이블)의 크기는 전체 모델 대비 미미하지만, 매우 작은 메모리 환경에서는 추가 최적화가 가능하다. 8. 결론 ZipServ는 LLM 가중치의 지수 비트가 낮은 엔트로피와 연속성을 활용해 고정 길이 비트맵 인코딩(TCA‑TBE)을 설계하고, 압축된 데이터를 Tensor Core 레지스터로 바로 복원하는 ZipGEMM 커널을 구현함으로써, 손실 없는 압축이 메모리 절감뿐 아니라 실제 추론 속도 가속으로도 이어질 수 있음을 입증한다. 이는 차세대 LLM 서비스에서 비용 효율적인 GPU 활용을 위한 중요한 전환점이 될 것으로 기대된다.