⏱️ [최신] LLVMhammer — 메모리 한 비트가 '컴파일러의 판단'을 뒤집으면, GPU는 멀쩡한데 청구서가 오른다

2026-10-24 · 카테고리: Research (최신 보안 논문 · Rowhammer · 컴파일러/LLM 서빙 · 하드웨어-소프트웨어 교차)

논문 정보

제목: Amplifying Latency and Cost of LLM Inference via Rowhammer Attack
저자: Mrityunjay Shukla, Shubham Roy, Biswabandan Panda (Indian Institute of Technology Bombay)
출처: arXiv:2610.30798 [cs.CR] · 2026년 10월 22일 제출 · 8 pages, 7 figures, 6 tables
분야: Rowhammer · Triton/MLIR-LLVM 컴파일러 · vLLM 추론 서빙 · 클라우드 공동 테넌시 · 가용성/비용 공격
인용: 2026-10-22 신규 공개 — 인용 집계 전 (OpenAlex/S2 레코드 미형성)
원문: arXiv:2610.30798 · HTML 전문 · 라이선스 CC BY-NC-SA 4.0

왜 이 논문인가 — "정확도는 그대로인데 느려졌다"는 가장 찾기 어려운 공격

LLM 서빙에서 보통 공격자는 모델 가중치를 노립니다. 비트를 뒤집어 답을 틀리게 만들거나, 백도어를 심거나, 컨텍스트를 빼냅니다. 그런 공격은 출력이 이상해지므로 결국 누군가 알아챕니다.

IIT Bombay의 Shukla·Roy·Panda가 10월 22일 올린 LLVMhammer는 정반대를 겨냥합니다. 뒤집는 대상은 모델도, GPU 메모리도 아닙니다. CPU DRAM에 올라와 있는 Triton 컴파일러 라이브러리(libtriton.so)입니다. 결과는 기능적으로 완전히 동일하지만 더 느린 GPU 커널 — 정확도 지표는 흔들리지 않고, TTFT·TPOT·처리량만 조금씩 나빠집니다.

논문의 한 줄 요약: "공격자는 코드를 바꾸지 않고, 컴파일러가 내리는 '최적화 결정'만 바꾼다." 그리고 Triton은 커널을 ~/.triton/cache에 저장하므로, 비트플립 한 번이 프로세스를 재시작해도 계속 남습니다.

위협 모델 — 같은 물리 호스트의 '평범한' 공동 테넌트

메커니즘 — 128비트 로드를 4개의 스칼라 로드로 바꾸는 '한 비트'

Triton은 Python DSL에서 시작해 MLIR → LLVM IR → PTX/AMDGCN → 기계어로 내려가는 파이프라인을 씁니다. 이 중 타깃 독립 단계(Triton IR·Triton-GPU IR 최적화)는 NVIDIA·AMD 백엔드가 공통으로 소비하는 표현입니다. 여기에 비트 하나가 뒤집히면, 벤더별 백엔드에 도달하기 전에 이미 잘못된 최적화 결정이 굳어집니다.

단계정상 커널비트플립 후
LLVM IRld.global.v4.b32 — 128비트 벡터 로드 1회ld.global.b32 × 4 — 스칼라 로드 4회
PTX128비트 코얼레스드 로드/스토어32비트 로드 4개, 메모리 트랜잭션 4배
SASSLDG.E.128스칼라 LDG 여러 개

Karnage — 118만 개 후보를 1,176개로 줄이는 자동화

libtriton.so에는 12만 개가 넘는 함수가 있고, 타깃 독립 구간만 봐도 후보 비트플립 위치가 100만 개를 넘습니다. 하나당 30초씩 추론을 돌리면 409일이 걸립니다. 저자들은 4단계 프루닝 도구 Karnage를 만들어 이 탐색을 현실적인 시간으로 줄입니다.

프루닝 단계NVIDIA 후보AMD 후보
Stage 1: 네임스페이스 필터 (타깃 독립 함수)1,183,0081,326,122
Stage 1: Intel PT 필터 (실행된 명령)38,76942,996
Stage 3: 오프라인 컴파일 출력 비교1,176810
Stage 4: 실제 GPU 프로파일링생존생존

실측 결과 — 4~16% 느려지고, 청구서는 최대 224만 달러

어디까지 믿어야 하나 — "Rowhammer로 실제로 뒤집었다"는 아직 아니다

이 논문은 정직하게도 자기 주장의 경계를 여러 곳에서 긋습니다. 실무자 관점에서 가장 중요한 구분은 공격의 두 조각이 분리돼 있다는 점입니다.

방어 — 정확도 모니터링이 아니라 '성능 기준선'을 감시해야 한다

교육용 PoC — 벡터가 스칼라로 쪼개질 때 '얼마나 손해인가'를 계산해 본다

아래 코드는 실제 하드웨어를 건드리지 않습니다. (1) 128비트 벡터 로드가 스칼라 로드로 바뀌었을 때의 메모리 트랜잭션 증가를 모델링하고, (2) 컴파일러 라이브러리·커널 캐시의 무결성 기준선을 검사하는 방어자 쪽 루틴을 흉내 냅니다. 목적은 "정확도는 그대로인데 느려지는" 변화를 어떤 지표로 잡을 수 있는가를 이해하는 것입니다.

import hashlib, os, statistics

# 1) 벡터 폭이 4 -> 1로 쪼개질 때의 상대 비용 모델 (교육용 근사)
#    트랜잭션 수는 4배, 대역폭 효율은 1/4. 실제 지연은 캐시/코얼레싱에 따라 달라진다.
def relative_cost(vec_width, elements=4096):
    transactions = elements // vec_width
    return transactions * vec_width / 32.0   # 32B 섹터 기준 정규화

base = relative_cost(4)
broken = relative_cost(1)
print(f"vectorized: {base:.1f}  scalarized: {broken:.1f}  "
      f"modeled slowdown: {broken/base:.2f}x")
# 실측 논문 값은 1.04~1.16x -- 모델은 상한(최악)에 가깝다.
# 즉 '트랜잭션 4배'가 곧 '지연 4배'는 아니다. 메모리 수준 병렬성이 흡수한다.

# 2) 방어자 루틴: 컴파일러 바이너리 + 커널 캐시의 기준선 해시
def manifest(paths):
    out = {}
    for p in paths:
        if os.path.isfile(p):
            with open(p, "rb") as f:
                out[p] = hashlib.sha256(f.read()).hexdigest()[:16]
    return out

def audit(baseline, actual):
    drift = []
    for p, h in actual.items():
        if baseline.get(p) not in (None, h):
            drift.append((p, baseline[p], h))
    return drift

# 예: 배포 시점에 baseline.json 을 만들어 두고, 매 세션 시작 시 비교한다.
baseline = {"libtriton.so": "a1b2c3d4e5f60718"}
actual = {"libtriton.so": "a1b2c3d4e5f60718"}   # 비트플립이 있으면 해시가 달라진다
print("integrity drift:", audit(baseline, actual))

# 3) 성능 기준선 감시: 같은 배치 크기에서의 지속적 지연 증가만 경보
def drifting(samples, baseline_ms, threshold=1.05, windows=3):
    hits = sum(1 for s in samples[-windows:] if s > baseline_ms * threshold)
    return hits == windows and len(samples) >= windows

print("slowdown alarm:", drifting([100, 101, 104, 106, 107], baseline_ms=100))
print("normal jitter :", drifting([101, 99, 103, 100, 102], baseline_ms=100))

그래서 무엇이 남는가

출처

목록으로 돌아가기