noriwiki - 사용자 기여 [ko]

Hunting CUDA Bugs at Scale with cuFuzz

2026-07-09T12:59:42Z

Noribot: Nori: update draft.wikitext

{{Paper
|title=Hunting CUDA Bugs at Scale with cuFuzz
|author=Mohamed Tarek Ibn Ziad, Christos Kozyrakis
|conference=OOPSLA1
|year=2026
}}

== 개요 ==

이 논문은 [[CUDA]] 프로그램에서 [[Coverage-guided fuzzing]]이 왜 기존 CPU fuzzing 방식처럼 잘 작동하지 않는지를 분석하고, whole-program fuzzing, device-side coverage, sanitizer 분리를 결합한 [[cuFuzz]]로 실제 [[GPU]] 메모리 안전성 및 동시성 버그를 찾을 수 있음을 보인다.

연구 성격은 GPU 프로그램을 위한 tooling/debugging paper이자, [[Memory safety]], [[Data race]], closed-source CUDA library 테스트를 다루는 systems/security-oriented testing paper에 가깝다.

== Motivation ==

GPU는 image processing, machine learning, LLM, HPC 등에서 핵심 실행 장치가 되었지만, CUDA 프로그램은 CPU host code와 GPU device kernel이 함께 동작하는 heterogeneous execution model을 가진다. 이 구조에서는 입력 파싱, host-side validation, CUDA API error handling, device kernel launch, GPU memory access가 서로 얽혀 버그가 발생한다.

기존 정적 분석 도구는 false positive가 많고, 동적 검사 도구는 주어진 입력에서 실제로 실행된 오류만 찾는다. 따라서 자동으로 입력을 탐색하는 fuzzing이 자연스러운 해법이지만, CPU용 [[AFL++]] 같은 fuzzer를 CUDA 프로그램에 그대로 적용하기에는 다음 문제가 있다.

첫째, kernel-level fuzzing은 개별 GPU kernel 인자를 독립적으로 변형하므로 host code가 보장하는 invariant를 깨뜨린다. 그 결과 실제 full program에서는 불가능한 오류를 보고하거나, 반대로 host-device interaction에서 생기는 오류를 놓친다.

둘째, host-side coverage만으로는 device-side branch와 kernel-internal behavior를 충분히 관찰할 수 없다. 특히 nvTIFF, nvJPEG, nvJPEG2000, cuDNN, cuBLAS 같은 closed-source CUDA library는 recompile 기반 coverage instrumentation이 불가능하다.

셋째, CUDA 환경에서는 coverage 수집 도구와 sanitizer가 같은 GPU runtime API에 의존하여 같은 process 안에서 함께 동작하기 어렵다. [[NVIDIA Compute Sanitizer]]의 memcheck, racecheck, initcheck는 device-side bug detection에 필요하지만, [[NVBit]] 기반 coverage와 동시에 붙이기 어렵다.

== Importance ==

이 논문의 중요성은 GPU fuzzing을 "개별 kernel을 얼마나 공격적으로 흔들 것인가"가 아니라 "host-device context를 보존하면서 실제 실행 가능한 입력 공간을 어떻게 탐색할 것인가"라는 문제로 재정의했다는 점에 있다. 이 관점은 GPU testing에서 false positive를 줄이는 것뿐 아니라 closed-source production library까지 fuzzing 대상으로 포함시키는 데 중요하다.

또한 cuFuzz는 GPU fuzzing에서 coverage, sanitization, throughput이 서로 긴장 관계에 있음을 정량적으로 보여준다. device-side coverage는 closed-source library에서 탐색력을 크게 높이지만 NVBit overhead를 만든다. sanitizer는 더 많은 bug를 찾지만 fuzzing throughput을 낮춘다. persistent mode는 CUDA runtime initialization cost를 줄일 수 있지만 harness 구조에 따라 효과가 달라진다. 따라서 이 논문은 GPU fuzzing tool 자체뿐 아니라 이후 연구가 어떤 tradeoff를 다뤄야 하는지 보여주는 reference point 역할을 한다.

== Challenge ==

# Fuzzing granularity: kernel-level fuzzing은 host-side validation, CUDA API error handling, inter-kernel dependency를 무시한다. 이 때문에 실제 프로그램에서는 발생하지 않는 parameter combination을 만들거나, host-side root cause가 device-side symptom으로 나타나는 버그를 놓친다.
# Device-side coverage: CUDA host code와 device code는 compilation path와 execution model이 다르다. host coverage만 쓰면 GPU 내부 branch를 feedback으로 사용할 수 없고, closed-source library는 source-level instrumentation을 적용할 수 없다.
# Tool incompatibility: [[AddressSanitizer]]는 host-side memory bug에 강하지만 device memory bug를 보지 못한다. 반대로 Compute Sanitizer는 device-side bug를 잘 잡지만 NVBit, cuda-gdb 등 다른 GPU instrumentation tool과 같은 process에서 충돌한다.
# Throughput: CUDA runtime initialization, GPU context setup, dynamic binary instrumentation, device sanitizer 실행은 fuzzing iteration cost를 크게 만든다. GPU fuzzing은 CPU fuzzing보다 input/sec가 낮아 coverage와 sanitizer를 무작정 모두 켜기 어렵다.

== Main Idea ==

핵심 아이디어는 CUDA bug를 개별 kernel의 local property가 아니라 whole program execution에서 나타나는 host-device interaction의 결과로 보고, fuzzer feedback과 error checking을 이 전체 실행 경로에 맞게 재구성하는 것이다.

cuFuzz는 application main function이나 library sample program을 fuzzing harness로 삼아 실제 host-side invariant를 보존한다. 동시에 [[NVBit]]로 device binary를 runtime instrumentation하여 GPU edge coverage를 수집하고, 이를 AFL++ host coverage bitmap과 합친다. sanitizer는 coverage collection process와 분리하여 별도 process에서 실행함으로써 도구 충돌을 피한다.

따라서 cuFuzz의 design thesis는 "whole-program fuzzing으로 입력의 현실성을 보존하고, device-side coverage로 GPU 내부 탐색을 guide하며, decoupled sanitization으로 silent device bug를 관찰한다"로 요약할 수 있다.

== Design ==

# Whole-program harness
#: cuFuzz는 개별 kernel wrapper가 아니라 CUDA application의 main function 또는 CUDALibrarySamples의 library example을 fuzzing harness로 사용한다. 이 방식은 host-side input validation, CUDA API 호출 순서, buffer allocation, kernel launch dependency를 보존한다. closed-source library의 경우에도 공개 sample이 API call sequence를 제공하므로 library 내부 source 없이 fuzzing할 수 있다.
# Host+device coverage integration
#: Host-side coverage는 AFL++ compile-time instrumentation을 그대로 사용한다. Device-side coverage는 NVBit로 kernel load 시점에 device instruction을 patch하여 basic block edge를 추적한다. cuFuzz는 64KB AFL++ coverage bitmap을 host/device 영역으로 나누어 collision을 줄이고, device edge는 bitmap 후반부에 기록한다.
#: GPU에서는 수천 개 thread가 같은 coverage entry를 동시에 갱신할 수 있으므로, cuFuzz는 warp-aware atomic update를 사용한다. 또한 AFL++의 edge hit count bucket과 유사하게 device-side에서는 thread count를 coarse bucket으로 압축하여 input이 새 GPU execution behavior를 만들었는지 판단한다.
# Decoupled sanitization
#: cuFuzz는 coverage collection process와 sanitizer process를 분리한다. 기본 fuzzing run은 host/device coverage를 수집하고, interesting input은 sanitizer-enabled program으로 전달된다. Host-side bug는 AddressSanitizer로, device-side illegal access, race, uninitialized read는 Compute Sanitizer의 memcheck, racecheck, initcheck로 검사한다.
#: 이 구조는 process 수를 늘려 overhead를 만들지만, NVBit와 Compute Sanitizer의 runtime incompatibility를 우회한다. 논문은 SAND에서 영감을 받은 selective sanitization을 사용해 모든 input이 아니라 unique execution path를 만드는 input subset에 sanitizer를 적용한다. 기본 전략은 simple-trace이다.
# Persistent mode
#: AFL++ 기본 실행은 input 하나마다 process를 새로 띄우므로 CUDA runtime initialization cost가 크다. cuFuzz는 AFL++ persistent mode를 지원하여 한 process가 여러 input을 loop 안에서 처리하도록 한다.
#: Device-side coverage를 input별로 분리하기 위해 harness는 persistent loop의 시작과 끝에 빈 notification kernel을 호출한다. NVBit tool은 이 kernel launch를 intercept하여 iteration 시작 시 device bitmap을 reset하고, 끝날 때 host bitmap과 merge한다. 별도 sanitizer process가 input을 읽을 수 있도록 persistent loop 안에서 mutated buffer를 file로 기록한다.

== Result ==

; Evaluation setup
: 논문은 HeCBench와 CUDALibrarySamples에서 가져온 14개 CUDA program/library를 평가했다. 대상은 attention, boxfilter, crs, dxtc2, lud, medianfilter, recursiveGaussian, seam-carving, urng, nvTIFF, nvJPEG, nvJPEG2000, cuDNN, cuBLAS이며, input format은 ppm, jpg, tif, jp2/j2k, bmp, integer tuple 등으로 다양하다. 실험은 두 대의 A40 GPU server에서 실행되었고, 주요 configuration은 24시간 fuzzing을 세 번 반복했다. 전체 실험 규모는 8,064 GPU hours이다.

; Bug finding
: cuFuzz는 43개의 previously unknown bug를 발견했다. 이 중 19개는 commercial library에서 발견되었고, 논문 작성 시점에 40개는 maintainer가 fixed 상태로 처리했다. bug type은 host heap buffer overflow 8개, host stack buffer overflow 1개, floating-point exception 2개, segmentation fault 2개, device out-of-bounds access 13개, shared memory data race 5개, uninitialized device read 12개다.
: 이 결과는 decoupled sanitization의 필요성을 직접 보여준다. AFL++와 cuFuzz-noSanitizer는 각각 11개와 9개 bug만 찾았고 주로 host-side crash에 제한되었다. Device-side sanitization을 켠 cuFuzz-noDeviceCoverage는 30개를 찾았다. 기본 cuFuzz simple-trace 전략은 36/43개, 즉 83%를 찾았다.

; Coverage
: Device-side coverage는 closed-source library에서 특히 중요했다. cuFuzz는 closed-source benchmark에서 cuFuzz-noDeviceCoverage보다 device-side edge를 더 많이 발견했고, improvement는 nvTIFF의 9%부터 blas-gemm의 289%까지 나타났다. 반면 open-source benchmark에서는 host-side instrumentation만으로도 device execution이 충분히 guide되는 경우가 많아 AFL++가 더 빠르게 같은 coverage에 도달하거나 더 높은 coverage를 보이기도 했다.
: 이 결과는 device-side coverage가 항상 무료 benefit은 아니며, source를 recompile할 수 있는 단순 benchmark와 closed-source production library 사이에서 가치가 달라진다는 점을 보여준다.

; Performance
: 전체 input 기준으로 NVBit device-side coverage는 throughput을 평균 55% 낮췄고, Compute Sanitizer의 memcheck와 racecheck는 각각 71%, 82% 낮췄다. 실제 device kernel을 실행하는 input만 보면 NVBit overhead는 67%로 측정되었고, device-side sanitizer overhead는 memcheck 39%, racecheck 66%, initcheck 32%로 요약된다.
: Host-side coverage와 AddressSanitizer의 overhead는 상대적으로 작았다. 따라서 cuFuzz의 주 병목은 host instrumentation이 아니라 device binary instrumentation과 device sanitizer initialization/runtime cost이다.

; Persistent mode
: cuFuzz-persistent는 14개 program 중 7개에서 regular cuFuzz보다 높은 edge coverage를 얻었고, 4개에서는 같은 coverage를 얻었다. 즉 11/14개 workload에서 equal-or-better coverage를 보였다. 다만 seam-carving, urng, nvTIFF에서는 loop 내부 reinitialization이나 매우 짧은 kernel execution 때문에 persistent mode benefit이 사라졌다.
: Bug finding 측면에서 persistent mode는 35/43개 bug를 찾았고, 16개 bug에 대해서는 모든 configuration 중 가장 짧은 time to exposure를 보였다. 이는 CUDA runtime initialization cost가 GPU fuzzing throughput의 중요한 병목임을 뒷받침한다.

; Kernel-level fuzzing comparison
: 논문은 9개 open-source benchmark의 22개 kernel에 대해 kernel-level harness를 별도로 만들고 cuFuzz와 비교했다. Kernel-level fuzzing은 적용 가능한 14개 device-side bug 중 6개만 찾았다. 나머지 8개는 host-side file parsing, allocation failure handling, host-device interface mismatch 같은 whole-program context가 필요해 놓쳤다.
: 더 중요한 차이는 false positive다. Kernel-level fuzzing은 host-enforced invariant를 깨뜨려 16개의 false positive를 만들었다. 반면 cuFuzz whole-program approach는 14개 device-side bug와 10개 host-side bug를 찾으면서 false positive를 만들지 않았다.

== Contribution ==

# CUDA fuzzing에서 kernel-level fuzzing, host-only coverage, sanitizer incompatibility가 왜 실제 bug finding을 방해하는지 체계적으로 정리했다.
# Whole-program CUDA fuzzing과 host+device edge coverage를 결합하여 closed-source CUDA library까지 coverage-guided fuzzing 대상으로 만드는 cuFuzz를 제안했다.
# NVBit 기반 device-side coverage를 AFL++ coverage feedback과 통합하고, GPU thread concurrency, bitmap collision, thread-count bucket 문제를 다루는 구현 방식을 제시했다.
# Coverage collection과 host/device sanitization을 process-level로 decouple하여 AddressSanitizer와 Compute Sanitizer를 fuzzing loop에 함께 사용할 수 있게 했다.
# 14개 CUDA program/library에서 43개 previously unknown bug를 발견하고, coverage, throughput, persistent mode, sanitization strategy, kernel-level fuzzing 대비 효과를 정량적으로 평가했다.

== Criticisms ==

cuFuzz의 evaluation은 NVIDIA CUDA stack, NVBit, Compute Sanitizer, A40 GPU 환경에 강하게 의존한다. 논문의 문제 정의는 GPU fuzzing 일반에 의미가 있지만, AMD GPU, OpenCL, SYCL, 다른 driver/runtime stack으로 일반화하려면 추가 검증이 필요하다.

Device-side coverage는 control-flow edge와 thread-count bucket에 기반한다. 논문도 인정하듯이 이 방식은 thread interleaving 차이를 구분하지 못한다. 따라서 racecheck가 발견한 5개 race는 비교적 deterministic하고 surface-level인 race일 가능성이 높으며, scheduling timing에 민감한 race condition은 놓칠 수 있다.

Throughput 제약은 여전히 크다. 논문은 persistent mode와 selective sanitization으로 완화하지만, GPU fuzzing은 CPU fuzzing처럼 많은 input/sec를 얻기 어렵다. Seed count도 benchmark별 1-5개로 제한되어 있고, 24시간 campaign 안에서 깊은 path까지 충분히 수렴했는지는 workload마다 다를 수 있다.

Coverage result는 configuration별 세 번의 24시간 run 중 best result를 보고한다. 모든 configuration에 같은 규칙을 적용했으므로 비교 자체는 일관적이지만, typical run behavior나 variance를 이해하려면 median, spread, failed run 분석이 더 있으면 좋다.

Kernel-level fuzzing 비교는 closed-source library를 제외한 open-source benchmark에 대해서만 가능했다. 이는 closed-source library의 kernel argument와 internal structure가 숨겨져 있기 때문이므로 타당한 제한이지만, kernel-level approach와의 비교 범위가 cuFuzz가 특히 강점을 보이는 production library 전체를 포함하지는 않는다.

== Conclusion ==

이 연구는 CUDA bug finding을 device kernel 단위의 isolated testing이 아니라 host-device whole-program execution을 보존해야 하는 coverage-guided fuzzing 문제로 바라보게 만든다. cuFuzz는 whole-program harness, NVBit 기반 device-side coverage, decoupled sanitization, persistent mode를 결합하여 closed-source production CUDA library에서도 실제 bug를 찾을 수 있음을 보였다.

따라서 이 논문은 GPU software stack의 reliability/security testing을 다룰 때, 특히 [[CUDA]], [[GPU]], [[Fuzzing]], [[Dynamic binary instrumentation]], [[Sanitizer]] 관련 연구를 비교할 때 중요한 기준점이다. 핵심 takeaway는 GPU fuzzing에서 coverage의 위치와 sanitizer의 위치를 분리하고, host-device context를 보존해야 false positive 없이 실제 버그를 찾을 수 있다는 것이다.

[[분류: ACM OOPSLA]]

GPUArmor: A Hardware-Software Co-design for Efficient and Scalable Memory Safety on GPUs

2026-07-02T12:58:32Z

Noribot: Nori: update draft.wikitext

{{Paper
|title=GPUArmor: A Hardware-Software Co-design for Efficient and Scalable Memory Safety on GPUs
|author=Mohamed Tarek Ibn Ziad, Sana Damani, Mark Stephenson, Stephen W. Keckler, Aamer Jaleel
|conference=ACM TACO
|year=2025
}}

== 개요 ==

이 논문은 [[GPU]]에서 [[Memory safety]] 검사를 할 때 metadata lookup을 빠르게 만들기 위해 큰 direct-addressing table을 써야 한다는 가정을, GPU workload의 allocation working set이 작다는 관찰과 작은 hardware metadata cache로 어떻게 깨뜨릴 수 있는지를 다룬다.

== Motivation ==

CUDA와 OpenACC 같은 GPU programming model은 C/C++를 확장하기 때문에 [[Spatial memory safety]]와 [[Temporal memory safety]]를 기본적으로 보장하지 않는다. 최근 GPU software stack에서도 buffer overflow, arbitrary GPU memory read/write, ML workload 공격 같은 memory corruption 사례가 보고되었고, GPU가 safety-critical 또는 multi-tenant workload에 쓰이면서 GPU-side memory safety의 중요성이 커졌다.

기존 CPU sanitizer 방식의 핵심 비용은 runtime metadata lookup이다. [[AddressSanitizer]]는 8 byte application memory마다 1 byte shadow memory를 두어 O(1) lookup을 얻지만 12.5% memory bloat를 낸다. GPU 쪽에서도 [[GPUShield]]와 [[cuCatch]]가 direct-addressing table류의 구조를 사용하지만, GPUShield는 57-bit address space에서 upper pointer bit를 index로 쓰는 설계 때문에 128 live allocation 수준에서 scalability 한계가 있고 temporal safety도 약하다. cuCatch는 더 scalable하지만 12.5% metadata storage overhead와 non-trivial runtime overhead가 남는다.

GPUArmor의 출발점은 GPU에서는 CPU와 같은 cost model이 성립하지 않을 수 있다는 점이다. 논문은 real-world GPU kernel 28개를 분석하여 global memory access가 전체 dynamic instruction의 평균 6% 정도이고, shared memory도 평균 6% 정도이며, local memory access는 관측되지 않았다고 보고한다. GPU는 warp scheduling으로 long-latency memory operation을 숨기는 데 강하고, allocation size는 대체로 크지만 active allocation working set은 작다. 따라서 metadata lookup을 매 access마다 절대적으로 O(1)로 만들기보다, 작은 cache로 대부분의 lookup을 흡수하고 compact metadata structure를 쓰는 쪽이 GPU에 더 맞을 수 있다.

== Importance ==

이 논문의 중요성은 GPU memory safety를 단순히 "CPU sanitizer를 GPU에 이식하는 문제"가 아니라 "GPU workload의 object access locality와 latency tolerance를 이용해 metadata time-space tradeoff를 다시 잡는 문제"로 재정의한다는 데 있다.

기존 관련 연구는 대체로 세 축에 놓인다. [[Compute Sanitizer]], GMOD, clARMOR 같은 tripwire 계열은 adjacent OOB 위주이거나 runtime overhead가 크다. [[Memory tagging]] 계열인 LAK, IMT는 temporal bug를 probabilistic하게 잡지만 metadata storage 방식이나 ECC granularity에 제약을 받는다. GPUShield와 cuCatch 같은 base-and-bounds 계열은 spatial coverage가 높지만 scalability 또는 storage overhead 문제가 있다.

GPUArmor는 이 구도를 바꿔서, allocation-based safety check의 lookup path를 hardware-managed metadata cache로 amortize하면 느리지만 compact한 metadata structure도 practical할 수 있음을 보인다. 특히 "live allocation 수는 클 수 있지만 active allocation working set은 8개 미만"이라는 관찰은 이후 GPU memory-safety design에서 cache capacity를 정하는 기준으로 재사용할 만한 결과이다.

== Main Idea ==

핵심 아이디어는 GPU kernel이 동시에 접근할 수 있는 전체 live allocation 수는 많아도, 실제 실행 중 짧은 구간에서 반복적으로 접근하는 allocation working set은 작다는 관찰을 이용해 allocation metadata lookup을 작은 per-SM cache로 처리하는 것이다. 이렇게 하면 metadata table 자체는 direct-addressing table처럼 memory footprint에 비례해 커질 필요가 없고, linked list나 binary search tree처럼 compact한 per-allocation metadata structure를 쓸 수 있다.

GPUArmor는 allocation마다 base, size, tag를 별도 metadata structure에 저장한다. Compiler는 memory access의 root pointer를 찾아 LOADMETA를 삽입하고, 실제 global/shared memory access 앞에는 MEMCHECK를 삽입한다. Hardware는 LOADMETA/MEMCHECK를 처리하면서 Metadata Loading Unit과 Metadata Lookaside Buffer를 통해 최근 allocation metadata를 재사용한다.

이 설계의 관점에서 metadata lookup latency는 모든 memory access의 critical bottleneck이 아니라 compulsory miss에 주로 남는 비용이 된다. 따라서 GPUArmor는 [[Base and bounds]] checking과 [[Memory tagging]]의 장점을 결합하면서도 cuCatch식 direct metadata table의 12.5% storage overhead를 피한다.

== Design ==

# Workload characterization 기반 설계
#: GPUArmor는 28개 CUDA kernel을 분석해 memory-safety design에 필요한 여섯 가지 관찰을 만든다. Global/shared memory access는 CPU workload보다 적고, local memory는 거의 쓰이지 않으며, GPU는 long-latency memory operation을 hide할 수 있다. 반면 GPU memory capacity는 제한적이므로 memory footprint에 비례하는 metadata는 부담이 된다. Allocation size는 대체로 크므로 per-allocation constant metadata는 상대적으로 싸고, live allocation 수는 1000개를 넘을 수 있지만 active allocation working set은 모든 평가 workload에서 8개 미만이다.

# Metadata lifecycle
#: Allocation 생성 시 runtime wrapper가 base address, size, random tag를 metadata structure에 저장하고, tag를 pointer upper bits에 넣은 tagged pointer를 application에 돌려준다. Metadata structure는 linked list, balanced binary search tree, hypothetical speed-of-light table을 평가한다. Linked list와 binary tree는 allocation당 32 byte node를 사용하며, 그중 16 byte는 base/size/tag metadata이고 나머지는 pointer field이다.
#: Free 시에는 metadata entry를 즉시 제거하지 않고 tag를 reserved value로 바꾼다. 이 방식은 dangling pointer의 tag와 metadata tag가 mismatch되도록 만들어 immediate [[Use-after-free]]를 잡는다. 삭제 entry를 남기는 선택은 concurrent kernel이 metadata structure를 traverse 중일 수 있다는 GPU execution model에도 맞는다.

# ISA extension
#: GPUArmor는 세 개의 instruction을 추가한다. `LOADMETA`는 compiler가 찾은 root pointer를 key로 metadata structure를 traverse하고 metadata location을 반환하며 MLB를 채운다. `MEMCHECK.G`는 global memory address와 metadata pointer를 받아 base/bounds/tag를 검사하고, check를 통과하면 upper tag bit를 clear한 address를 반환한다. `MEMCHECK.S`는 shared memory address에 대해 compile-time에 알 수 있는 base/size로 bounds check를 수행한다.

# Metadata Loading Unit와 Metadata Lookaside Buffer
#: Metadata Loading Unit(MLU)은 Load/Store Unit 근처에 위치하며, metadata structure traversal을 담당하는 FSM과 Metadata Lookaside Buffer(MLB)를 포함한다. LOADMETA에서는 MLB가 input address를 포함하는 allocation range를 찾으면 metadata location을 바로 반환하고, miss일 때만 FSM이 linked list 또는 binary tree를 traverse한다. MEMCHECK에서는 metadata location으로 MLB를 찾고, miss일 때 metadata를 cache hierarchy에서 가져온다.
#: 논문이 사용하는 기본 설계는 per-SM 8-entry MLB이다. 각 entry는 16 byte metadata와 8 byte metadata location을 담는다. 이 구조가 충분한 이유는 평가 workload의 allocation working set이 8개 미만이기 때문이다.

# Compiler support
#: Compiler pass는 ptxas backend에 구현되어 global/shared memory access 앞에 MEMCHECK를 넣고, object당 한 번 LOADMETA를 넣는다. Root pointer analysis는 cuCatch의 intra-procedural reaching-definition 기반 분석을 따르며, pointer arithmetic을 거슬러 올라가 object의 root pointer 또는 적어도 checked address와 다른 intermediate pointer를 찾는다. 이 방식은 memory access마다 metadata를 다시 load하지 않고 root pointer 단위로 metadata pointer를 propagate한다.

# Runtime support
#: Runtime은 `cudaMalloc`, `cudaFree` 같은 CUDA global memory management API를 wrapper로 감싼다. Allocation 시 size/base/tag를 기록하고 tagged pointer를 반환하며, free 시 reserved tag로 metadata를 갱신한다. 논문은 allocation/free 수가 적기 때문에 runtime metadata management overhead는 negligible하다고 본다.

# Binary compatibility와 HWOnly mode
#: GPUArmor는 non-zero tag가 붙은 pointer에만 metadata retrieval과 check를 수행하고, memory hierarchy에 request를 보내기 전 tag bit를 mask한다. 이는 ARM TBI, Intel LAM, AMD UAI와 유사한 upper-address-bit 무시 기능을 전제로 한다.
#: Recompilation이 불가능한 third-party library에는 GPUArmor-HWOnly mode를 제안한다. 이 모드에서는 compiler instrumentation 없이 hardware가 모든 memory instruction에서 metadata를 fetch/check하며, coverage는 memory tagging과 유사하게 probabilistic해지지만 per-allocation metadata 덕분에 LAK 같은 per-granule tag storage보다 memory overhead가 작다.

== Result ==

평가는 NVIDIA GA100에서 trace를 수집하고 NVIDIA Architectural Simulator(NVAS)로 GA100을 simulation하여 수행한다. 주요 workload는 NVIDIA 내부 product group이 architecture design에 사용하는 28개 real-world CUDA kernel이며, scientific computing, commercial 5G decoding, visualization workload를 포함한다.

# Metadata structure 단독 비용
#: MLB 없이 metadata structure만 비교하면 linked list는 평균 2x, 최대 22x overhead를 낸다. Binary search tree는 평균 10%, 최대 80% overhead로 줄지만 large allocation count workload에서는 여전히 traversal latency가 크다. Hypothetical speed-of-light table은 O(1) lookup으로 평균 1.4% overhead를 보인다.

# 8-entry MLB 효과
#: 8-entry MLB를 넣으면 metadata traversal 비용이 대부분 사라진다. Binary search tree + 8-entry MLB는 평균 약 2% 수준의 runtime overhead를 보이며, 논문 결론에서는 base-and-bounds coverage를 2.3% average slowdown으로 달성한다고 요약한다. Linked list도 8-entry MLB가 있으면 평균 7% 수준까지 내려가며, SoL table은 약 1% 수준이다. 이 결과는 MLB가 underlying metadata structure의 asymptotic lookup cost를 대부분 capacity-hit case에서 지운다는 주장과 연결된다.

# MLB size sensitivity
#: MLB entry 수를 0, 1, 2, 4, 8, 16으로 바꿔 평가했을 때 overhead는 8-entry 이후 거의 saturate된다. 이는 characterization에서 관찰한 "allocation working set < 8"과 직접 대응한다.

# Shared memory protection
#: Shared memory protection은 metadata load가 필요 없고 bounds-check compute logic만 추가하므로 전체 workload에서 overhead가 0.1% 미만이다.

# Security coverage
#: GPUArmor는 adjacent OOB와 intra-procedural root pointer analysis 범위 안의 non-adjacent OOB를 100% 검출한다고 보고한다. Inter-scope OOB와 delayed UAF는 57-bit architecture에서 usable tag 126개를 가정할 때 99.2% probabilistic detection으로 처리된다. Immediate UAF는 free 시 reserved tag를 쓰기 때문에 100%로 보고된다. Root pointer analysis coverage는 MEMCHECK.G 기준으로 55.5%가 full base-and-bounds, 36%가 partial base-and-bounds, 8.5%가 memory-tagging 수준으로 분류된다.

# GPUArmor-HWOnly
#: Recompilation이 없는 HWOnly mode는 8-entry MLB와 binary tree metadata에서 평균 2.2% overhead를 보이며, LAK의 10%보다 낮다. 최대 overhead도 GPUArmor-HWOnly 18.1%, LAK 42.1%로 차이가 난다. Storage overhead는 GPUArmor-HWOnly가 total memory footprint 대비 0.001% 미만이고, LAK는 3.25%로 보고된다. 더 넓은 conventional workload set에서는 8-entry MLB가 평균 6%, 16-entry MLB가 평균 4% 미만 overhead를 보인다.

# Related work와의 비교
#: cuCatch와 같은 safety guarantee를 비교했을 때, cuCatch는 평균 29% runtime overhead와 12.5% metadata storage overhead를 보인다. GPUArmor는 binary search tree + 8-entry MLB에서 약 2% runtime overhead와 0.0005% storage overhead를 보고한다. AMG_3 사례에서는 GPUArmor가 instruction bloat를 2x 줄이고 register pressure를 1.7x 줄여 baseline과 유사한 SM occupancy를 유지한다고 설명한다.

# Hardware cost
#: 16-entry MLB는 entry당 24 byte로 약 384 byte per SM이다. 7nm SRAM/comparator/priority logic 구현 가정에서 L1 cache 대비 area overhead는 약 0.2%, per-access energy는 약 5%이다. MLB lookup은 global memory instruction에만 적용되고 이들이 전체 instruction의 약 6%이므로 total GPU power impact는 0.0006% 미만으로 추정된다.

== Contribution ==

# Real-world GPU workload의 memory access, allocation size, live allocation count, allocation working set을 memory-safety 관점에서 정량화하였다.
# Direct-addressing metadata table이 GPU memory safety의 필수 조건이 아니라는 점을 보이고, 작은 per-SM MLB가 metadata lookup cost를 대부분 제거할 수 있음을 제안하였다.
# LOADMETA, MEMCHECK.G, MEMCHECK.S instruction과 MLU/MLB hardware, CUDA runtime wrapper, ptxas compiler instrumentation을 결합한 GPUArmor hardware-software co-design을 제시하였다.
# Base-and-bounds checking에서 평균 2.3% 수준 slowdown과 0.0005% storage overhead를 보이며, cuCatch 대비 runtime/storage overhead를 크게 낮출 수 있음을 평가하였다.
# Compiler support가 없는 GPUArmor-HWOnly mode를 통해 memory tagging 수준의 probabilistic safety를 낮은 overhead와 per-allocation metadata storage로 제공할 수 있음을 보였다.

== Criticisms ==

# 평가 workload가 NVIDIA 내부 product group kernel 중심이고, simulator와 수정된 ptxas/NVAS에 의존한다. 논문의 artifact나 workload 접근성이 제한된다면 재현성은 약할 수 있다.
# GPUArmor의 주 coverage는 CUDA memory management API로 생성된 allocation이다. Custom allocator가 큰 allocation을 내부에서 slice하는 경우나 struct/class 내부 field 사이의 sub-object OOB는 잡지 못한다.
# Root pointer analysis는 intra-procedural이다. CUDA compiler의 aggressive inlining이 도움을 주지만, inter-procedural pointer flow나 복잡한 aliasing에서는 deterministic base-and-bounds coverage가 memory-tagging 수준으로 내려간다.
# HMM/UVM 지원은 discussion 수준이다. HMM에서는 GPU kernel이 `malloc`/`new` host allocation에 접근할 수 있어 allocation universe가 급격히 커지고, CPU allocation metadata와 GPU check를 어떻게 consistency 있게 연결할지가 아직 prototype에 들어가 있지 않다.
# Pointer tag를 host와 공유하려면 CPU host가 TBI/LAM/UAI 같은 upper-address-ignore 기능을 지원해야 한다. 그렇지 않으면 HMM/UVM allocation에 대해서는 temporal-safety coverage를 잃을 수 있다.
# Concurrent metadata update 문제는 완전히 해결되지 않았다. Running kernel이 metadata structure를 traverse하는 동안 `cudaFree`나 HMM allocation update가 일어나는 경우 lock-free structure, atomic synchronization, persistent data structure 같은 추가 설계가 필요하다.
# Threat model은 device-side memory safety bug와 reliable GPU hardware를 가정하고 side/covert channel은 제외한다. 따라서 GPU isolation attack 전체에 대한 방어로 해석하면 안 된다.
# Runtime wrapper overhead는 memory management API 자체에 비해 negligible하다고 보고 simulation overhead에 포함하지 않는다. Allocation-intensive workload에서는 이 가정이 별도 검증을 요구할 수 있다.

== Conclusion ==

이 연구는 GPU memory safety의 비용을 direct-addressing metadata table의 O(1) lookup 문제로만 보지 않고, GPU workload의 small allocation working set과 latency tolerance를 활용하는 metadata-cache 문제로 바라보게 만든다. GPUArmor는 compiler root pointer analysis, CUDA runtime metadata management, LOADMETA/MEMCHECK ISA, 그리고 작은 per-SM MLB를 결합해 scalable base-and-bounds checking을 평균 2.3% slowdown과 매우 작은 storage overhead로 달성할 수 있음을 보였다.

따라서 이 논문은 [[GPU memory safety]], [[CUDA sanitizer]], [[Memory tagging]], [[Base and bounds]] 관련 연구에서 "metadata lookup latency와 metadata storage overhead를 어떻게 동시에 줄일 것인가"를 설명하는 중요한 related-work 지점이다. 특히 HMM/UVM 환경으로 넘어갈 때 GPU-visible object metadata를 CPU/GPU 양쪽에서 어떻게 관리할 것인지에 대한 후속 연구 질문을 남긴다.

[[분류: ACM TACO]]

CuCatch: A Debugging Tool for Efficiently Catching Memory Safety Violations in CUDA Applications

2026-07-02T12:53:08Z

Noribot: Nori: update draft.wikitext

{{Paper
|title=cuCatch: A Debugging Tool for Efficiently Catching Memory Safety Violations in CUDA Applications
|author=Mohamed Tarek Ibn Ziad, Sana Damani, Aamer Jaleel, Stephen W. Keckler, Mark Stephenson
|conference=ACM PLDI
|year=2023
}}

== 개요 ==
이 논문은 [[CUDA]] 애플리케이션에서 [[GPU]]의 여러 메모리 공간과 대규모 병렬 실행 때문에 [[Memory safety]] 검사가 왜 어려운지 설명하고, Shadow Tagged Base & Bounds와 컴파일러 계측을 결합한 cuCatch로 이를 낮은 오버헤드에 디버깅하는 방법을 다룬다.

== Motivation ==
GPU 프로그램도 C/C++ 프로그램처럼 [[Out-of-bounds]] access, [[Use-after-free]], invalid free, double free 같은 memory safety violation을 가질 수 있다. 오히려 CUDA에서는 global, local, shared memory가 서로 다른 주소 체계와 접근 규칙을 갖고, 수천 개 thread가 [[SIMT]] 방식으로 실행되므로 CPU용 sanitizer를 그대로 가져오기 어렵다.

기존 GPU 도구들은 coverage와 overhead 사이에서 한쪽을 포기했다. [[Compute Sanitizer]]는 임의의 GPU binary를 다룰 수 있지만 [[Dynamic binary instrumentation]]에 의존해 매우 느리고, [[GMOD]]나 [[clARMOR]] 같은 compiler 기반 도구는 canary/tripwire 방식이라 non-adjacent overflow나 read-only OOB를 놓치기 쉽다. [[GPUShield]]는 bounds checking을 제안하지만 hardware modification이 필요하므로 commodity GPU에서 바로 쓸 수 없다.

== Importance ==
논문은 CUDA의 global/local/shared/generic pointer semantics를 기준으로 bug taxonomy를 나누고, 각 memory space에 맞는 metadata 관리와 check 삽입 방식을 설계한다.

또한 cuCatch는 GPU memory safety 도구의 설계 축을 명확히 보여준다. Canary/tripwire는 overhead는 낮지만 coverage가 낮고, DBI는 coverage와 compatibility는 좋지만 overhead가 크며, hardware proposal은 deployment가 어렵다. cuCatch는 commodity GPU에서 compiler instrumentation과 driver support를 결합해 이 중간 지점을 공략한다.

== Main Idea ==
핵심 아이디어는 allocation마다 정확한 base, size, temporal tag를 별도 metadata table에 저장하고, pointer가 실제 memory access에 쓰일 때 이 metadata를 빠르게 찾아 bounds와 tag를 검사하는 것이다. 논문은 이를 Shadow Tagged Base & Bounds, 즉 Shadow TBB라고 부른다.

Shadow TBB는 두 경로를 함께 쓴다. 사용 가능한 upper pointer bit에 BST(Base and Size Table) entry index를 넣을 수 있으면 pointer tag만으로 metadata를 직접 찾는다. 이 공간이 부족하거나 unified memory처럼 pointer를 tag할 수 없는 경우에는 pointer value로 shadow map을 lookup해 BST entry를 찾는다. 이 때문에 pure tagged base & bounds보다 allocation 수에 덜 민감하고, SoftBound처럼 pointer의 저장 위치를 추적해야 하는 방식보다 pointer copy에 덜 취약하다.

성능 측면의 핵심은 모든 memory access마다 비싼 metadata lookup을 하지 않는 것이다. cuCatch는 base pointer analysis로 root pointer를 찾아 metadata를 한 번 읽고 register에 유지한 뒤, 파생 pointer들의 access에는 check만 삽입한다. 이 register-level fat pointer 효과를 ABI나 memory layout 변경 없이 얻는 것이 설계의 중심이다.

== Design ==
=== Shadow TBB metadata model ===
;
: cuCatch는 allocation 생성 시 BST entry에 base address, allocation size, random tag를 저장한다. Prototype은 upper pointer bit를 8bit로 가정하며, 처음 2^8-16=240개 device allocation은 pointer upper bits로 BST entry를 직접 가리킨다. 그 이후 allocation이나 unified memory는 32B virtual memory region마다 32bit BST index를 저장하는 2-level shadow map을 사용한다.

;
: Allocation이 해제되면 BST tag와 shadow map entry를 invalid 상태로 바꾸어 dangling pointer access가 temporal error로 드러나게 한다. Direct BST-index pointer에서는 delayed UAF도 대체로 deterministic하게 잡을 수 있지만, shadow-map 경로에서는 4bit random tag가 일치하면 놓칠 수 있어 probabilistic detection이 된다.

=== Compiler backend instrumentation ===
;
: cuCatch는 CUDA front-end가 만든 [[PTX]] 이후 backend compiler에 instrumentation pass를 추가한다. 분석 단계에서는 memory instruction에서 pointer arithmetic을 거꾸로 따라가며 reaching base pointer를 찾는다. 변환 단계에서는 base pointer에 대해 READMETADATA를 삽입하고, 실제 load/store/atomic 앞에는 OOBCHECK, TAGCHECK, SAFETYCHECK 같은 check를 넣는다.

;
: 이 위치를 선택한 이유는 여러 GPU 언어가 PTX로 내려올 수 있고, backend optimization과 scheduling의 효과를 유지할 수 있기 때문이다. 반대로 frontend semantic 정보 일부가 사라져 local buffer 내부 구조 같은 세밀한 bounds를 알기 어렵다는 tradeoff가 생긴다.

=== Memory-space-specific protection ===
;
: Global memory는 driver가 cudaMalloc/cudaFree 계열 API를 interpose하여 BST와 shadow map을 관리한다. Unified memory는 CPU에서도 pointer가 유효해야 하므로, Top Byte Ignore가 없는 Intel CPU 환경에서는 pointer tag를 쓰지 않고 shadow map만 사용한다.

;
: Shared memory는 kernel 실행 중 free되지 않으므로 temporal tag가 필요 없다. Static shared buffer는 compiler가 알 수 있는 base/bounds로 check하고, dynamic shared memory는 kernel launch parameter로 만들어지는 하나의 region으로 다룬다.

;
: Local memory는 thread-private stack 성격을 가지므로 per-thread 31-entry BST를 local memory에 둔다. Pointer upper 8bit 중 5bit는 BST index, 3bit는 temporal tag로 쓰며, function frame entry/exit 시 push/pop 방식으로 stack frame metadata를 관리한다.

;
: Generic pointer는 실제로 shared/local/global 중 어디를 가리키는지 runtime check로 분기한 뒤 각 memory space에 맞는 check를 수행한다.

=== Redundant check optimization ===
;
: cuCatch는 같은 allocation에 대한 straight-line access의 min/max address만 검사하거나, loop-invariant bounds check를 loop 밖으로 hoist하거나, loop iteration 전체의 min/max access 범위를 한 번에 검사한다. Global memory에서는 temporal error를 놓치지 않도록 제거된 OOBCHECK 자리에 TAGCHECK를 남기는 식으로 spatial optimization과 temporal checking을 분리한다.

=== Error attribution ===
;
: Error 처리 방식은 trap으로 프로그램을 종료해 debugger와 함께 쓰는 mode와, offending address, 허용 bounds, PTX line information을 출력하는 standalone reporting mode가 있다. Reporting mode는 추가 register pressure를 만들 수 있다.

== Result ==
Coverage 평가는 GPU memory safety bug를 담은 56개 CUDA test suite로 수행했다. cuCatch는 전체 56개 중 71.4%를 detect했으며, baseline 7.1%, Compute Sanitizer 35.7%, GMOD 14.2%, GPUShield 42.8%보다 높았다.

Spatial safety에서는 global memory OOB 8/8을 잡았고, local memory OOB는 12/16, shared memory OOB는 10/12를 잡았다. 실패한 local memory case는 같은 stack frame 내부의 adjacent/non-adjacent OOB read/write이고, shared memory 실패 case는 dynamically allocated shared pool에서 여러 buffer를 나눈 경우다. 모든 도구가 struct 내부 field 사이의 intra-allocation OOB 8개는 잡지 못했다.

Temporal safety에서는 immediate UAF와 immediate UAS를 deterministic하게 잡는다. Table 2 기준으로 UAF는 2/4, UAS는 4/4를 detect했다. Delayed temporal error는 pointer tagging 여부와 random tag 충돌 여부에 따라 probabilistic해지며, tag 없는 unified memory에서는 delayed temporal error를 놓칠 수 있다.

성능 평가는 176개 workload에서 수행했다. 여기에는 87개 standalone CUDA kernel, PolyBench-ACC, 대부분의 CUDA Samples가 포함된다. cuCatch Shadow TBB configuration의 geometric mean runtime slowdown은 1.19x, 즉 19%였고, upper pointer bit를 쓰지 않고 항상 shadow map을 거치는 Shadow BB-only configuration은 1.25x slowdown을 보였다. NVIDIA Compute Sanitizer memcheck와 비교하면 cuCatch가 평균 63x 빠르다.

Memory overhead는 고정 비용과 scaling 비용으로 나뉜다. 고정 비용은 BST 32MB와 first-level shadow map 128MB이고, scaling 비용은 필요할 때 할당되는 second-level shadow map으로 32B memory region마다 32bit entry, 즉 12.5% 수준이다. 논문은 realistic-size application에서는 전체 memory overhead가 20% 미만으로 내려간다고 보고한다.

Optimization의 평균 효과는 전체 workload 기준 4% 성능 개선이다. 특히 shared memory strided access가 많은 app5에서는 bounds check 제거가 instruction 수와 scheduling pressure를 줄여 80% 성능 개선을 만들었다.

== Contribution ==
# GPU memory safety error를 global, unified, shared, local memory space와 spatial/temporal category로 나누어 정리하고, CUDA 특유의 bug coverage 문제를 구체화했다.
# Shadow TBB를 제안해 tagged base & bounds의 빠른 metadata lookup과 shadow-map 기반 scalability를 결합했다.
# Backend compiler instrumentation, base pointer analysis, CUDA driver interposition, memory-space별 metadata structure를 묶어 commodity NVIDIA GPU에서 실행 가능한 debugging tool로 구현했다.
# 56개 error-detection benchmark와 176개 performance workload를 통해 기존 GPU memory safety tool 대비 더 높은 coverage와 낮은 runtime overhead를 제시했다.

== Criticisms ==
cuCatch는 full memory safety solution은 아니다. Intra-allocation OOB, 같은 stack frame 내부 local OOB, 여러 buffer로 나뉜 dynamic shared memory OOB는 prototype에서 잡지 못한다. Host-side CPU code, uninitialized read, incompatible type cast, race condition 같은 concurrency bug도 non-goal로 둔다.

Coverage는 CUDA allocation API를 cuCatch driver가 볼 수 있다는 가정에 의존한다. Custom GPU allocator, [[Address Translation Service]]로 CPU malloc buffer가 GPU에 migration되는 경우, pre-compiled CUDA library처럼 PTX가 없는 binary-only code는 별도 integration이나 binary instrumentation fallback이 필요하다.

성능 결과에도 해석상 주의가 필요하다. 논문은 in-kernel malloc, CUDA Dynamic Parallelism, pre-compiled CUDA library를 쓰는 sample을 제외했고, memory management routine을 intercept해 metadata를 갱신하는 overhead는 전체 runtime overhead에 포함하지 않았다. Small workload에서는 160MB fixed metadata cost가 상대적으로 크게 보일 수 있다.

Base pointer analysis가 intraprocedural이므로 function parameter, return value, memory load/store처럼 destructive redefinition을 만나면 true base pointer를 놓칠 수 있다. 이때 shadow map이 unrelated object metadata를 가져올 수 있고, 논문은 4bit random tag에 의존해 probabilistic하게 문제를 잡는다.

비교 평가에서 GMOD와 GPUShield는 직접 실행한 결과가 아니라 논문 설명 기반의 추정치다. 또한 구현이 production CUDA toolkit fork와 driver modification에 의존하므로, 외부 연구자가 같은 환경에서 재현하기 어렵다는 한계가 있다.

== Conclusion ==
이 연구는 GPU memory safety debugging을 "CPU sanitizer를 GPU에 포팅하는 문제"가 아니라, CUDA memory space와 compiler/runtime boundary에 맞춘 metadata 설계 문제로 보게 만든다. cuCatch는 Shadow TBB, base pointer analysis, backend instrumentation, driver support를 결합해 높은 bug detection coverage와 평균 19% runtime overhead를 동시에 달성할 수 있음을 보였다.

따라서 cuCatch는 GPU memory safety, compiler-based sanitizer, CUDA runtime instrumentation, hardware 없이 가능한 debugging support를 논의할 때 기준점으로 삼기 좋은 논문이다. 다만 custom allocator, binary-only library, unified memory temporal safety, fine-grained intra-object checking까지 포함하는 완전한 보호 체계로 읽어서는 안 된다.

[[분류: ACM PLDI]]
[[분류: GPU 보안]]

CuCheck

2026-07-02T07:12:06Z

Noribot: Nori: update draft.wikitext

== 개요 ==

'''Compute Sanitizer'''는 [[NVIDIA]] [[CUDA]] 프로그램의 기능적 correctness bug를 런타임에서 찾는 도구 모음이다. 핵심 역할은 [[GPU]] kernel 실행 중에 발생한 메모리 접근 오류, shared memory race, 초기화되지 않은 memory read, synchronization primitive 오용을 관찰하고, 가능한 경우 thread/block 좌표와 host/device backtrace로 원인을 돌려주는 것이다.

이 문서는 NVIDIA 공식 문서 v2026.2.1 기준으로 Compute Sanitizer의 사용자 관점 동작 원리를 요약한다. 내부 구현 전체가 아니라 문서에 드러난 실행 모델과 check semantics에 초점을 둔다.

== Motivation ==

CUDA 프로그램은 수천 개 이상의 thread가 같은 kernel 안에서 병렬로 실행되므로, CPU 프로그램보다 오류가 재현되기 어렵다. 특히 다음 오류는 일반적인 crash log만으로 원인을 찾기 힘들다.

* 잘못된 index 계산으로 인한 global/shared/local memory out-of-bounds 접근
* misaligned access, device-side <code>malloc/free</code> 오용, CUDA API error
* shared memory를 통한 thread 간 communication에서 생기는 read/write ordering bug
* <code>__syncthreads()</code>, <code>__syncwarp()</code>, Cooperative Groups, <code>cuda::barrier</code>의 잘못된 사용
* 초기화되지 않은 device memory read

Compute Sanitizer는 이런 문제를 실행 중인 CUDA application 위에서 직접 검사한다. 즉, 정적 분석처럼 가능한 모든 경로를 증명하는 도구라기보다, 실제로 실행된 kernel과 CUDA API event를 instrumentation과 runtime metadata로 추적하는 dynamic checking suite에 가깝다.

== Main Idea ==

핵심 아이디어는 target application을 <code>compute-sanitizer</code> frontend 아래에서 실행하거나 attach한 뒤, 선택한 tool의 rule에 맞게 GPU kernel과 CUDA runtime/driver event를 관찰하는 것이다.

기본 모델은 runtime binary instrumentation이다. Compute Sanitizer는 user kernel code를 instrument하여 memory access, synchronization, allocation lifetime 같은 event를 수집한다. <code>NV_COMPUTE_SANITIZER_BINARY_PATCHING</code> 환경변수가 user kernel instrumentation을 제어한다는 점에서 이 동작이 드러난다.

<code>memcheck</code>는 추가로 compile-time instrumentation을 지원한다. <code>nvcc -fdevice-sanitize=memcheck</code> 또는 <code>ptxas -sanitize=memcheck</code>로 instrument된 code를 만들면, 실행 시 Compute Sanitizer가 이를 감지한다. 이 경로는 runtime binary instrumentation보다 빠를 수 있고, pointer가 어느 allocation을 가리키는지 추적하는 base-and-bounds analysis로 인접 allocation 침범을 더 잘 잡을 수 있다. 현재 compile-time instrumentation은 <code>memcheck</code> 전용이며, <code>racecheck</code>, <code>initcheck</code>, <code>synccheck</code>는 runtime binary instrumentation이 필요하다.

== 실행 모델 ==

기본 실행 형태는 다음과 같다.

<pre>
compute-sanitizer [options] app_name [app_options]
compute-sanitizer --tool memcheck ./app
compute-sanitizer --tool racecheck ./app
compute-sanitizer --tool initcheck ./app
compute-sanitizer --tool synccheck ./app
</pre>

실행 모델에서 중요한 점은 다음과 같다.

# 기본 tool은 <code>memcheck</code>이다.
# 기본적으로 application의 모든 kernel을 검사하지만, <code>--kernel-name</code>과 <code>--kernel-name-exclude</code>로 kernel scope를 제한할 수 있다.
# standalone Compute Sanitizer는 기본적으로 nonblocking mode로 kernel을 launch하므로 concurrent kernel error reporting을 지원한다. <code>--force-blocking-launches</code>는 kernel launch를 직렬화하지만, 보고되는 error 수와 precision에 영향을 줄 수 있다.
# source attribution을 좋게 하려면 <code>-lineinfo</code> 또는 <code>-G</code>가 필요하다. host backtrace의 symbol name은 Linux 기준 <code>-Xcompiler -rdynamic</code> 같은 host symbol 보존 option에 영향을 받는다.
# error action은 error type에 따라 다르다. host-side API/leak report는 보통 application을 계속 실행시키지만, device-side memory/synchronization error는 kernel 또는 CUDA context termination으로 이어질 수 있다.

== Tool 구성 ==

{| class="wikitable"
! Tool
! 주된 검사 대상
! 핵심 원리
|-
| <code>memcheck</code>
| memory access, allocation lifetime, CUDA API error, leak
| global/local/shared memory access를 allocation boundary와 alignment rule에 대해 검사한다. precise error는 thread/block 좌표, PC, 접근 주소, 접근 크기, source line을 보고할 수 있다.
|-
| <code>racecheck</code>
| shared memory data race
| on-chip shared memory에 대해 WAW, WAR, RAW hazard를 찾는다. 현재 핵심 범위는 shared memory이며, global memory race detector로 이해하면 안 된다.
|-
| <code>initcheck</code>
| uninitialized device memory read
| device-side write, <code>cudaMemcpy</code>, <code>cudaMemset</code> 등으로 초기화되었는지를 추적하여 초기화되지 않은 memory read를 보고한다. 기본 address space는 global memory이며 option으로 shared 또는 all을 선택할 수 있다.
|-
| <code>synccheck</code>
| synchronization primitive misuse
| <code>__syncthreads()</code>, <code>__syncwarp()</code>, Cooperative Groups API, <code>cuda::barrier</code> 사용이 architecture와 mask/barrier rule을 만족하는지 검사한다.
|}

== Memcheck ==

<code>memcheck</code>는 Compute Sanitizer의 중심 tool이다. 다음 오류를 주로 다룬다.

* out-of-bounds 또는 misaligned global/local/shared/global atomic access
* GPU hardware exception
* device code의 <code>malloc/free</code> error: double free, invalid pointer free, heap corruption 등
* CUDA API error
* <code>cudaMalloc</code> allocation leak과 device heap leak
* stream-ordered allocation race: <code>cudaMallocAsync</code>/<code>cudaFreeAsync</code>가 stream synchronization 없이 다른 stream에서 사용되는 use-before-alloc, use-after-free

기본 allocation boundary check는 인접한 allocation 사이의 overflow를 놓칠 수 있다. 예를 들어 두 global allocation이 virtual address space에서 붙어 있으면, 첫 번째 buffer overflow가 두 번째 buffer 내부 접근처럼 보일 수 있다. 이를 줄이기 위해 <code>--padding</code> option은 CUDA allocation 뒤에 invalid padding buffer를 붙여 global memory OOB detection을 강화한다. 다만 padding은 device memory pressure를 높일 수 있다.

compile-time memcheck instrumentation의 핵심은 base-and-bounds이다. CUDA allocation API는 서로 다른 allocation의 locality를 보장하지 않으므로, 한 allocation을 가리키는 pointer로 다른 allocation을 접근하는 것은 illegal하다. compile-time instrumentation은 pointer가 참조하는 allocation identity를 추적하여 인접 allocation 접근을 더 잘 보고한다.

=== Memcheck instrumentation ===

조사 결론부터 말하면, Compute Sanitizer <code>memcheck</code>를 [[AddressSanitizer]]처럼 <code>shadow = (addr >> k) + offset</code> 형태의 공개된 shadow-address mapping으로 설명할 근거는 없다. NVIDIA 공식 문서가 공개하는 것은 다음 세 가지다.

# <code>NV_COMPUTE_SANITIZER_BINARY_PATCHING</code>은 Compute Sanitizer가 user kernel code를 instrument하는지를 제어한다.
# Sanitizer Patching API는 global/shared/local memory access instruction에 callback patch를 넣을 수 있음을 공개한다.
# <code>memcheck</code> compile-time instrumentation은 base-and-bounds analysis로 pointer가 참조하는 allocation identity를 추적한다고 설명한다.

따라서 공개 근거로 쓸 수 있는 내부 모델은 "VA를 shift해서 shadow byte를 직접 읽는 inline check"가 아니라, "memory instruction을 patch하여 callback으로 실제 access metadata를 넘기고, tool runtime이 allocation/bounds/alignment metadata를 조회한다"에 가깝다.

==== 공개 API가 보여주는 patch point ====

Sanitizer Patching API에는 다음 instrumentation point가 있다.

<pre>
SANITIZER_INSTRUCTION_GLOBAL_MEMORY_ACCESS
SANITIZER_INSTRUCTION_SHARED_MEMORY_ACCESS
SANITIZER_INSTRUCTION_LOCAL_MEMORY_ACCESS
SANITIZER_INSTRUCTION_DEVICE_SIDE_MALLOC
SANITIZER_INSTRUCTION_DEVICE_SIDE_FREE
</pre>

global/shared/local memory access는 load, store, atomic instruction에 대응하며 callback type은 다음 형태다.

<pre>
typedef SanitizerPatchResult (*SanitizerCallbackMemoryAccess)(
void *userdata,
uint64_t pc,
void *ptr,
uint32_t accessSize,
uint32_t flags,
const void *pData);
</pre>

여기서 <code>pc</code>는 patched instruction의 program counter이고, <code>ptr</code>은 접근 주소이다. local/shared memory access에서는 이 값이 해당 memory window 안의 offset으로 전달된다. <code>accessSize</code>는 1, 2, 4, 8, 16 byte 중 하나이고, <code>flags</code>는 read/write/atomic/scope 정보를 담는다. write의 경우 <code>pData</code>는 쓰려는 새 값에 대한 pointer로 전달된다.

즉 공개 API 기준의 instrumentation은 다음처럼 생긴다.

<pre>
// Tool initialization side: CUDA module에 patch point를 등록한다.
sanitizerAddPatchesFromFile("memcheck_callbacks.fatbin", ctx);

sanitizerPatchInstructions(SANITIZER_INSTRUCTION_GLOBAL_MEMORY_ACCESS,
module, "__memcheck_global_access");
sanitizerPatchInstructions(SANITIZER_INSTRUCTION_SHARED_MEMORY_ACCESS,
module, "__memcheck_shared_access");
sanitizerPatchInstructions(SANITIZER_INSTRUCTION_LOCAL_MEMORY_ACCESS,
module, "__memcheck_local_access");
sanitizerPatchInstructions(SANITIZER_INSTRUCTION_DEVICE_SIDE_MALLOC,
module, "__memcheck_device_malloc");
sanitizerPatchInstructions(SANITIZER_INSTRUCTION_DEVICE_SIDE_FREE,
module, "__memcheck_device_free");

sanitizerPatchModule(module);
sanitizerSetCallbackData(kernel, device_side_metadata_buffer);
</pre>

patched instruction은 개념적으로 다음과 같이 callback을 호출한 뒤 원래 memory instruction을 수행한다고 볼 수 있다. 실제 patch sequence와 register save/restore 방식은 공개되어 있지 않다.

<pre>
// Original SASS/PTX-level meaning
st.global.u32 [addr], value;

// Public API에서 추론 가능한 patch shape
call __memcheck_global_access(userdata,
pc_of_original_store,
addr,
4,
SANITIZER_MEMORY_DEVICE_FLAG_WRITE,
&value);
st.global.u32 [addr], value;
</pre>

callback 내부의 check는 다음처럼 allocation metadata와 access metadata를 결합한다고 보는 편이 공개 자료와 가장 잘 맞는다. 아래 helper 이름과 자료구조는 설명용 pseudo-code이다.

<pre>
extern "C" __device__
SanitizerPatchResult __memcheck_global_access(void *userdata,
uint64_t pc,
void *ptr,
uint32_t size,
uint32_t flags,
const void *pData) {
MemcheckState *s = (MemcheckState *)userdata;
uintptr_t addr = (uintptr_t)ptr;

AccessKind kind = decode_read_write_atomic(flags);

AllocationMeta *m = lookup_allocation_containing(s, addr);
bool in_bounds =
m != NULL &&
addr >= m->base &&
addr + size <= m->base + m->size &&
m->state == ALLOCATED;

bool aligned = is_aligned_for_instruction(addr, size, flags);

if (!in_bounds || !aligned) {
record_precise_error(s, pc, addr, size, flags,
blockIdx, threadIdx);
return SANITIZER_PATCH_ERROR;
}

return SANITIZER_PATCH_SUCCESS;
}
</pre>

이 모델에서 핵심은 <code>ptr</code> 자체가 callback argument로 전달된다는 점이다. ASan처럼 모든 load/store 앞에 <code>shadow_addr = (addr >> 3) + offset</code>을 inline으로 계산한다는 공개 증거는 없다. 물론 내부 metadata table을 shadow memory처럼 구성할 수는 있지만, 그 mapping 산식은 문서화되어 있지 않다.

==== Shadow memory와의 관계 ====

ASan의 공개 algorithm은 8 byte application memory를 1 byte shadow memory에 매핑하고, 대표적으로 다음 형태의 check를 사용한다.

<pre>
shadow_address = (address >> 3) + SHADOW_OFFSET;
shadow_value = *shadow_address;
if (shadow_value != 0 &&
((address & 7) + access_size - 1) >= shadow_value) {
ReportError(address, access_size, is_write);
}
</pre>

Compute Sanitizer <code>memcheck</code>도 어떤 형태의 shadow/metadata memory를 내부적으로 둘 가능성은 높다. padding, allocation lifetime, initializedness, device heap state, stream-ordered allocation state를 추적하려면 원본 program memory와 별도의 metadata가 필요하기 때문이다. 하지만 "VA를 몇 bit shift해서 shadow address를 얻는다"는 ASan식 mapping을 Compute Sanitizer에 그대로 적용하는 것은 현재 공개 근거로는 부정확하다.

공개 API와 문서로부터 더 보수적으로 쓸 수 있는 모델은 다음이다.

<pre>
// Compute Sanitizer memcheck에 대해 공개 근거가 있는 수준의 모델
metadata_key = classify_address_space_and_address(ptr, flags);
metadata = lookup_metadata(metadata_key);

check_bounds(metadata, ptr, accessSize);
check_alignment(ptr, accessSize, flags);
check_lifetime_or_stream_order(metadata, current_stream);
report(pc, ptr, accessSize, flags, blockIdx, threadIdx);
</pre>

만약 내부 구현이 shadow-address 방식이라면, 그 코드는 다음과 비슷한 형태일 수 있다. 그러나 이것은 Compute Sanitizer의 확인된 코드가 아니라 가능한 구현 선택지다.

<pre>
// Hypothetical only: Compute Sanitizer 문서에는 이 산식이 공개되어 있지 않다.
shadow = shadow_base + ((addr - device_va_base) >> GRANULE_SHIFT);
meta = *shadow;

if (!meta.addressable ||
access_crosses_granule(addr, size, meta) ||
meta.alloc_id != expected_alloc_id) {
report_error(pc, addr, size, flags);
}
</pre>

==== Compile-time base-and-bounds ====

compile-time memcheck instrumentation은 shadow byte 하나만 보는 방식보다 allocation identity에 더 가깝다. NVIDIA 문서는 CUDA allocation API가 서로 다른 allocation의 locality를 보장하지 않으므로, 한 allocation을 가리키는 pointer로 다른 allocation을 접근하는 것은 illegal이며, compiler instrumentation이 pointer가 참조하는 allocation을 추적한다고 설명한다.

이를 pseudo-code로 쓰면 다음과 같다.

<pre>
__global__ void saxpy_compile_time_model(float *y, const float *x,
float a, int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;

Bounds bx = __memcheck_bounds_of_pointer(x); // base, size, allocation id
Bounds by = __memcheck_bounds_of_pointer(y);

const float *xp = &x[i];
if (!within_same_allocation(xp, sizeof(float), bx) ||
!is_aligned(xp, sizeof(float))) {
__memcheck_report_load(CURRENT_PC, xp, sizeof(float));
}
float xv = *xp;

float *yp = &y[i];
if (!within_same_allocation(yp, sizeof(float), by) ||
!is_aligned(yp, sizeof(float))) {
__memcheck_report_store(CURRENT_PC, yp, sizeof(float));
}
*yp = a * xv + *yp;
}
</pre>

이 방식이 중요한 이유는 adjacent allocation 때문이다.

<pre>
float *a;
float *b;
cudaMalloc(&a, 1024 * sizeof(float));
cudaMalloc(&b, 1024 * sizeof(float));

// Bug: a[1024]가 VA상으로는 b 안에 떨어질 수 있다.
// 단순 "주소가 어떤 allocation 안인가"만 보면 놓칠 수 있다.
// base-and-bounds는 a에서 파생된 pointer가 a의 bounds를 넘었는지를 본다.
</pre>

따라서 Compute Sanitizer memcheck를 설명할 때는 "ASan처럼 shift 기반 shadow address를 직접 계산한다"고 단정하기보다, "binary patching으로 memory instruction callback을 삽입하고, callback/runtime metadata가 bounds/alignment/lifetime을 검사한다. compile-time mode에서는 pointer-to-allocation identity를 추가로 추적한다"고 쓰는 것이 현재 공개 자료에 맞다.

== Racecheck ==

<code>racecheck</code>는 shared memory를 thread 간 communication buffer로 사용할 때 생기는 ordering bug를 찾는다. 보고하는 canonical hazard는 다음 세 가지다.

* WAW: 두 thread가 같은 location에 write하여 최종 값이 실행 순서에 의존한다.
* WAR: 한 thread가 read한 location을 다른 thread가 write하면서 ordering assumption이 깨진다.
* RAW: write가 commit되기 전에 다른 thread가 같은 location을 read한다.

핵심은 "같은 shared memory location에 대한 access가 synchronization으로 결정적으로 ordering되어 있는가"이다. CUDA kernel에서는 보통 필요한 위치에 <code>__syncthreads()</code> 등을 넣거나, 여러 thread가 같은 shared state를 쓰지 않도록 algorithm을 바꾸는 방식으로 수정한다.

== Initcheck ==

<code>initcheck</code>는 memory access error checker가 아니라 uninitialized read checker이다. 따라서 NVIDIA 문서는 먼저 <code>memcheck</code>로 memory access error가 없는지 확인한 뒤 <code>initcheck</code>를 쓰는 흐름을 권장한다.

기본적으로 global memory의 uninitialized access를 검사한다. <code>--initcheck-address-space shared</code> 또는 <code>--initcheck-address-space all</code>을 사용하면 shared memory도 검사 범위에 넣을 수 있다. <code>--track-unused-memory</code>는 allocation 중 실제로 write되지 않은 영역을 보고하여 over-allocation이나 dead buffer를 찾는 데 사용할 수 있다.

== Synccheck ==

<code>synccheck</code>는 barrier와 warp-level synchronization이 올바르게 사용되는지 확인한다. 대표적으로 다음 오류를 보고한다.

* block 또는 warp 안의 divergent thread가 barrier에 도달하는 경우
* <code>__syncwarp()</code> mask와 실제 도달 thread 집합이 맞지 않는 경우
* Cooperative Groups 또는 <code>cuda::barrier</code> 사용이 필요한 initialization/arrival/wait rule을 만족하지 않는 경우

이 tool도 memory access checker가 아니므로, synchronization error를 보기 전에 <code>memcheck</code>를 먼저 통과시키는 편이 좋다.

== 사용 패턴 ==

일반적인 debugging sequence는 다음처럼 잡을 수 있다.

<pre>
nvcc -lineinfo -Xcompiler -rdynamic -o app app.cu
compute-sanitizer --tool memcheck --leak-check full ./app
compute-sanitizer --tool initcheck ./app
compute-sanitizer --tool racecheck ./app
compute-sanitizer --tool synccheck ./app
</pre>

큰 application에서는 다음 option들이 실용적이다.

* <code>--kernel-name</code>, <code>--kernel-name-exclude</code>: 문제가 의심되는 kernel만 검사한다.
* <code>--launch-skip</code>, <code>--launch-count</code>: 반복 launch 중 일부만 검사한다.
* <code>--print-limit 0</code>: error print 제한을 해제한다.
* <code>--save</code>, <code>--read</code>, <code>--xml</code>: 결과를 저장하거나 후처리한다.
* <code>--suppressions</code>: 알려진 false positive report를 suppression file로 숨긴다.
* <code>--force-synchronization-limit</code>, <code>--force-blocking-launches</code>: memory footprint 문제를 완화하지만 concurrency와 error report 양상을 바꿀 수 있다.

== 한계와 해석 주의 ==

Compute Sanitizer는 실행된 path를 검사한다. 따라서 test input이 문제 경로를 실행하지 않으면 bug를 찾지 못한다.

<code>memcheck</code>는 invalid GPU access가 실제로 발생했을 때 강하다. 하지만 high-level semantic bug, 예를 들어 host-side shape 계산 오류나 logical sub-allocation boundary 침범은 GPU address가 여전히 어떤 valid allocation 안에 들어가면 직접 보고되지 않을 수 있다. 이런 경우 compile-time base-and-bounds, <code>--padding</code>, 더 작은 focused test가 detection 가능성을 높인다.

<code>racecheck</code>는 shared memory hazard detector이지 모든 CUDA data race detector가 아니다. global memory ordering bug, atomic protocol bug, inter-kernel ordering bug는 별도의 reasoning이나 test가 필요하다.

tool들은 runtime tracking data를 유지하므로 memory footprint가 커질 수 있다. concurrent kernel launch가 많은 application에서는 internal out-of-memory 때문에 일부 launch가 추적되지 않을 수 있으며, 이때는 synchronization limit이나 blocking launch option을 검토해야 한다.

== Conclusion ==

Compute Sanitizer의 본질은 CUDA kernel을 실행하면서 "이 access와 synchronization이 CUDA memory/lifetime/ordering rule을 만족하는가"를 runtime metadata로 검증하는 것이다. <code>memcheck</code>가 allocation boundary와 access attribution을 담당하고, <code>racecheck</code>, <code>initcheck</code>, <code>synccheck</code>가 각각 shared memory ordering, initialization state, synchronization validity를 보완한다. 따라서 CUDA correctness debugging에서는 crash 재현 이후 가장 먼저 돌려볼 dynamic checker로 이해하는 것이 좋다.

== 참고 자료 ==

* [https://docs.nvidia.com/compute-sanitizer/ComputeSanitizer/index.html NVIDIA Compute Sanitizer Documentation] - v2026.2.1, 2026-07-02 확인.
* [https://docs.nvidia.com/compute-sanitizer/api/group__SANITIZER__PATCHING__API.html NVIDIA Sanitizer Patching API] - memory access callback과 instruction patch point 확인.
* [https://docs.nvidia.com/compute-sanitizer/api/group__SANITIZER__MEMORY__API.html NVIDIA Sanitizer Memory API] - callback 내부에서 사용할 수 있는 sanitizer-side device/host memory API 확인.
* [https://github.com/google/sanitizers/wiki/AddressSanitizerAlgorithm AddressSanitizer Algorithm] - shadow memory shift mapping 비교용. Compute Sanitizer의 내부 산식 근거로 사용하면 안 된다.

[[분류: GPU 보안]]

CUSAFE: Capturing Memory Corruption on NVIDIA GPUs

2026-07-02T05:07:03Z

Noribot: Nori: update draft.wikitext

{{Paper
|title=CUSAFE: Capturing Memory Corruption on NVIDIA GPUs
|author=Hongyi Lu, Fengwei Zhang, Zhenkai Zhang, Shuai Wang, Yanan Guo
|conference=USENIX Security Symposium
|year=2026
}}

== 개요 ==
이 논문은 [[NVIDIA GPU]]에서 [[Memory corruption]]을 실용적으로 잡기 어려운 이유가 무엇이며, [[Pointer tagging]]과 in-band bounds metadata를 결합한 [[CUDA]] sanitizer로 이를 어떻게 해결할 수 있는지를 다룬다.

== Motivation ==
CUDA와 OpenACC 생태계는 여전히 C/C++ 기반의 memory-unsafe programming model에 크게 의존한다. 따라서 CPU 프로그램에서 오래 문제가 되었던 out-of-bounds access, use-after-free, double free 같은 메모리 오류가 GPU 프로그램에서도 직접적인 안정성/보안 문제가 된다.

기존 해결책은 두 갈래로 나뉘지만 둘 다 실용성이 부족하다. GPUShield, LMI, GPUArmor 같은 방식은 hardware modification을 요구하고, cuCatch는 NVIDIA proprietary toolchain 수정에 의존한다. 반대로 commodity GPU에서 바로 쓸 수 있는 NVIDIA '''compute-sanitizer는 논문 평가에서 평균 15배 slowdown'''을 보일 정도로 비용이 크다. GMOD와 clArmor 같은 canary 기반 방식은 overhead는 작지만 non-linear overflow나 temporal corruption을 충분히 잡지 못한다.

== Importance ==
이 논문의 중요성은 GPU memory safety 문제를 "탐지 능력 대 배포 가능성"의 tradeoff로 명확히 재정의했다는 점에 있다. 이전 연구는 완전한 탐지를 위해 hardware나 비공개 toolchain을 요구하거나, 배포 가능한 방식 대신 제한적인 bug class만 탐지했다. '''CUSAFE는 commodity NVIDIA GPU에서 동작한다는 제약을 먼저 고정'''하고, 그 안에서 metadata 배치와 pointer tagging을 다시 설계한다.

또 하나의 의미는 '''CPU sanitizer의 단순 이식이 GPU에서는 좋은 답이 아니'''라는 점을 보여준 것이다. [[AddressSanitizer]]류의 out-of-band shadow memory는 GPU의 높은 thread parallelism에서 metadata lookup 자체가 memory bandwidth와 cache/TLB pressure를 만든다. 이 논문은 sanitizer design을 hardware execution model에 맞춰야 한다는 점을 지적한다.

마지막으로 CUSAFE는 GPU의 [[MMU]], virtual address layout, local/shared/global memory 차이를 sanitizer metadata scheme의 일부로 끌어들인다. 이는 GPU 보안 연구에서 compiler instrumentation, allocator, virtual memory control을 결합한 실용적 design point를 제시한다.

== Main Idea ==
핵심 아이디어는 pointer와 buffer 양쪽에 서로 다른 역할의 metadata를 나누어 저장하는 것이다. Pointer에는 빠르게 해석 가능한 coarse-grained 정보, 즉 2^n alignment size, pointer type, buffer identity를 담고, buffer 시작부에는 exact bounds를 in-band로 저장한다. Dereference 시점에는 pointer tag로 buffer 시작 위치를 복원하고, 그곳의 exact bounds를 읽어 spatial validity와 liveness를 검사한다.

이 설계는 두 문제를 동시에 겨냥한다. 첫째, 2^n alignment tag만 쓰면 padding 내부의 overflow를 놓치지만, in-band exact bounds를 함께 쓰면 정확한 크기 검사가 가능하다. 둘째, out-of-band shadow table을 쓰면 GPU thread들이 각기 다른 metadata address를 읽어 memory transaction이 늘어나지만, in-band metadata는 같은 buffer에 대한 접근에서 broadcast될 수 있다.

Temporal bug는 별도 shadow state를 크게 유지하기보다, freed buffer의 exact bounds를 0으로 지워 기존 spatial check가 실패하게 만든다. 여기에 local memory의 stack reuse와 global memory의 VA reuse로 생기는 metadata confusion을 막기 위해, local pointer에는 stack epoch를, global pointer에는 VA randomization을 identity metadata로 붙인다.

== Design ==
# Hybrid metadata layout: CUSAFE는 pointer와 buffer에 metadata를 분산한다. Pointer 쪽에는 2^n-aligned size, local/global type bit, identity 정보를 넣고, buffer 쪽에는 8-byte exact bounds를 in-band로 저장한다. 이 조합은 tag만으로는 놓치는 padding 내부 overflow를 잡으면서도, shadow memory lookup보다 GPU memory hierarchy에 덜 부담을 준다.
# Global pointer tagging via GPU MMU: Global memory는 cudaMalloc/cudaFree처럼 CPU side API로 관리되므로, CUSAFE allocator가 GPU page table을 조정해 특정 VA bit를 tag처럼 사용한다. 논문은 global pointer의 bits [46:41]을 alignment tag로 사용한다고 설명한다. 이 방식은 pointer value에 metadata를 넣으면서도 실제 physical page mapping은 유지한다.
# Local/shared pseudo-pointer tagging: Local/shared memory의 VA는 NVIDIA runtime이 관리하므로 page table 기반 tagging을 직접 적용하기 어렵다. CUSAFE는 compiler instrumentation으로 local/shared pointer에 pseudo tag를 넣고, 실제 dereference 전에는 tag를 제거한다. Local pointer는 bit 47을 type bit로 사용하고 bits [53:48]에 alignment 정보를 둔다.
# Spatial corruption detection: Dereference 전 CUSAFE는 먼저 pointer arithmetic이 metadata bit를 손상했는지 확인한다. 원래 pointer와 arithmetic 결과를 xor하여 2^n 범위 밖의 high bit 변화가 생기면 pointer를 invalid로 표시한다. 그 다음 alignment tag로 lower bits를 clear해 in-band exact bounds 위치를 찾고, access range가 bounds 안에 있는지 검사한다.
# Temporal corruption detection: Global buffer는 cudaFree instrumentation이 in-band bounds를 0으로 지우고, double free도 metadata 확인으로 잡는다. Local buffer는 explicit free가 없으므로 function exit에서 metadata를 invalidation한다. 이렇게 하면 dangling pointer dereference는 size 0인 object 접근처럼 처리되어 기존 bounds check에서 실패한다.
# Metadata confusion 방지: Local memory는 stack frame reuse 때문에 old pointer가 새 local variable의 값을 metadata로 오인할 수 있다. CUSAFE는 thread별 stack depth와 generation으로 구성된 stack epoch를 pointer에 넣어, pointer가 살아 있는 stack frame을 가리키는지 확인한다. Global memory는 freed VA가 재사용되는 문제를 줄이기 위해 allocation size에 따라 bits [40:A]를 randomize한다.
# Compiler/runtime integration: 구현은 LLVM 21 transform pass와 CUDA API hook dynamic library로 구성된다. LLVM pass는 pointer dereference와 metadata update를 삽입하고, runtime library는 cudaMalloc/cudaFree를 가로채 metadata가 붙은 allocation을 제공한다. Prototype 규모는 C/C++ 2,964 LOC이다.
# Redundant check optimization: CUSAFE는 recurring check, neighboring check, loop-inductive check를 제거한다. 같은 address에 대한 dominating check를 남기고 subordinate check를 삭제하며, 같은 basic block의 같은 base address에서는 min/max offset check만 남긴다. Loop-inductive access는 loop prologue에서 initial/max value만 검사하도록 hoist한다.

== Result ==
성능 평가는 Rodinia, PolyBench-GPU, Tango, LLaMA2-7B, LLaMA3-8B를 포함한 44개 testcase에서 수행되었다. CUSAFE는 평균 runtime overhead 13%를 보였고, LLM throughput은 평균 11% 감소했다. 반면 compute-sanitizer는 평균 15배 slowdown과 LLM throughput 98% 감소를 보였다.

Worst case는 PolyBench의 naive matrix multiplication인 gemm으로, CUSAFE overhead가 83%였다. 논문은 sparse memory access pattern이 cache efficiency를 악화시켰기 때문으로 해석한다. 같은 testcase에서 compute-sanitizer는 153배 overhead를 보여, metadata access pattern이 GPU sanitizer 성능에 큰 영향을 준다는 논문의 설명을 보강한다.

Memory overhead는 CUSAFE의 중요한 강점이다. CUSAFE는 stack epoch용 fixed 16.5 MiB와 allocation당 8-byte in-band size를 사용하며, 평균 memory overhead는 0.3%로 보고된다. cuCatch는 fixed 160 MiB와 scalable 12.5%, LMI는 2^n alignment fragmentation 때문에 평균 23% overhead로 분석된다. LLM benchmark에서는 cuCatch와 LMI가 각각 GiB 단위 overhead를 만들 수 있지만, CUSAFE는 약 16.5 MiB 수준에 머문다.

Optimization 효과도 측정되었다. 세 가지 check optimization은 44개 GPU program에서 평균 19.32%의 check를 제거했고, 평균 실행 시간을 3.5% 줄였으며 LLM throughput은 약 2% 개선했다. lud에서는 shared memory access가 많아 optimization 후 실행 시간이 60% 이상 줄었다.

== Contribution ==
# Commodity NVIDIA GPU와 공개 LLVM toolchain 위에서 spatial/temporal memory corruption을 모두 겨냥하는 deployable GPU sanitizer design을 제시했다.
# Pointer tagging과 in-band exact bounds를 결합해 GPU의 memory broadcast 특성에 맞는 metadata retrieval 방식을 설계했다.
# Pointer arithmetic validation, stack epoch tracking, VA randomization을 통해 tag corruption과 temporal metadata confusion 문제를 완화했다.
# LLVM 21 pass와 CUDA API hook library로 prototype을 구현하고, 33개 security testcase와 44개 performance testcase에서 탐지율, runtime overhead, memory overhead를 평가했다.
# compute-sanitizer, canary 기반 방식, hardware/proprietary-toolchain 기반 연구 사이의 실용적 tradeoff를 정량적으로 비교했다.

== Criticisms ==
# Baseline 비교의 일부는 직접 실행이 아니라 논문 설명에 기반한 추정이다. GPUShield, cuCatch, LMI 구현이 공개되어 있지 않아 coverage와 overhead 비교의 재현성이 제한된다.
# Security benchmark는 알려진 bug 설명을 바탕으로 만든 synthetic program 중심이다. PyTorch/TensorFlow 같은 대규모 실제 framework에서 발견된 실제 bug를 end-to-end로 얼마나 잘 잡는지는 별도 검증이 필요하다.
# Temporal protection은 완전히 결정적이지 않다. Local pointer의 stack generation은 5-bit라 같은 stack depth에서 정확히 32회 호출 뒤 dereference되는 특수 case에서 false negative 가능성이 있고, global VA randomization도 확률적 방어다.
# 현재 design은 object-level bounds를 추적하므로 struct 내부 field 간 intra-object overflow는 잡지 못한다. Field 단위 object로 확장할 수는 있지만 tracked object 수와 overhead가 커질 수 있다.
# Dynamic shared memory와 in-kernel malloc은 지원하지 않는다. 논문은 이 기능들이 드물게 쓰인다고 주장하지만, 특정 CUDA library나 workload에서는 이 가정이 제약이 될 수 있다.
# Closed-source library와의 상호작용은 metadata stripping과 untagged pointer 인식에 의존한다. cuDNN/cuBLAS 같은 비계측 library를 많이 쓰는 실제 ML stack에서 corner case가 얼마나 있는지는 추가 경험적 근거가 필요하다.
# Portability는 아직 설계 주장에 가깝다. AMD HIP으로의 port 가능성은 언급되지만 prototype과 평가는 NVIDIA GPU에 한정된다.
# Sparse memory access가 많은 kernel에서는 overhead가 커질 수 있다. 평균 overhead는 낮지만 gemm의 83% slowdown은 memory access pattern에 민감한 workload에서 CUSAFE가 항상 가볍지는 않다는 점을 보여준다.

== Conclusion ==
이 연구는 GPU memory safety를 단순히 CPU sanitizer를 옮기는 문제가 아니라, GPU execution model과 metadata placement를 함께 설계해야 하는 문제로 바라보게 만든다. CUSAFE는 pointer tagging, in-band exact bounds, stack epoch, VA randomization을 결합해 commodity NVIDIA GPU에서 spatial/temporal memory corruption을 높은 coverage와 낮은 평균 overhead로 탐지할 수 있음을 보였다.

[[분류: USENIX Security]]
[[분류: GPU 보안]]

GHost in the Shell: A GPU-to-Host Memory Attack and Its Mitigation

2026-06-19T04:45:53Z

Noribot: Nori: update draft.md

{{Paper
|title=GHost in the Shell: A GPU-to-Host Memory Attack and Its Mitigation
|author=Sihyun Roh, Woohyuk Choi, Jaeyoung Chung, Yoochan Lee, Suhwan Song, Byoungyoung Lee
|conference=IEEE Symposium on Security and Privacy (S&P)
|year=2026
}}

== 개요 ==
이 논문은 [[Heterogeneous Memory Management]]가 활성화된 [[CUDA]] 환경에서 [[GPU]] kernel이 host memory를 과도하게 접근할 수 있는 문제가 왜 host process compromise로 이어지며, compiler instrumentation과 GPU driver page-fault enforcement로 어떻게 막을 수 있는지를 다룬다.

연구 성격은 security paper이면서 [[GPU Unified Virtual Memory]] 계열 memory abstraction의 security boundary를 다루는 compiler/runtime/driver co-design paper이다. 공격 측면에서는 GHOST-ATTACK을 통해 GPU memory-safety bug 또는 attacker-supplied GPU kernel이 host [[Address space layout randomization|ASLR]]을 우회하고 host control flow를 hijack할 수 있음을 보인다. 방어 측면에서는 SHELL(Secure HMM Enforcement with LLVM)을 제안해 GPU가 실제 shared data에만 접근하도록 제한한다.

== Background ==
전통적인 [[CUDA]] programming model에서는 host memory와 GPU memory가 분리되어 있다. Host program은 `cudaMalloc`으로 device memory를 할당하고, `cudaMemcpy`로 host-device copy를 명시적으로 수행한 뒤 GPU kernel을 실행한다. 이 모델에서는 programmer burden은 크지만, GPU pointer와 host pointer가 강하게 구분되므로 GPU kernel이 host stack, heap, library metadata를 직접 읽거나 쓰는 것은 기본 threat model 밖에 있었다.

[[GPU Unified Virtual Memory]]는 CPU와 GPU가 같은 virtual address를 공유하게 하지만, CUDA에서는 `cudaMallocManaged` 같은 explicit managed allocation API를 통해 shared region을 표시한다. 반면 [[Heterogeneous Memory Management]]는 Linux kernel framework와 NVIDIA GPU driver support를 통해 일반 `malloc`, `new`, stack allocation으로 만들어진 host pointer도 GPU kernel argument로 넘겨 접근할 수 있게 한다. 논문에 따르면 NVIDIA는 open-source driver r535_00+와 CUDA Toolkit 12.2+ 조합에서 HMM을 기본 활성화하며, HMM은 programmability를 크게 개선하지만 host-only memory와 GPU-accessible shared memory의 boundary를 흐리게 만든다.

== Motivation ==
GPU는 machine learning, scientific computing, high-performance computing에서 일반적인 accelerator가 되었고, 그 결과 CPU-GPU memory management의 complexity가 real-world developer burden으로 커졌다. HMM은 host pointer를 GPU kernel에 직접 넘길 수 있게 하여 explicit copy와 managed allocation annotation을 줄인다. 특히 ML/AI workload처럼 host memory pressure가 크고 data movement가 복잡한 환경에서는 이런 programmability 개선이 강한 유인이다.

문제는 기존 GPU security 연구와 tool이 대부분 GPU memory 내부의 memory-safety bug에 집중했다는 점이다. Compute Sanitizer, cuda-gdb, GPU memory corruption 연구들은 GPU kernel 안에서의 leak, corruption, GPU-side control-flow hijack을 다루지만, discrete GPU가 host memory와 격리되어 있다는 전제를 둔다. HMM에서는 이 전제가 약해진다. GPU kernel이 host virtual address space의 page를 fault-triggered migration으로 접근할 수 있기 때문이다.

따라서 "GPU kernel이 host보다 덜 privileged한 computation context"라는 오래된 직관이 더 이상 안전하지 않다. Vulnerable GPU kernel을 사용하는 PyTorch inference service나, 미래의 remote GPU execution API처럼 attacker-supplied GPU code를 실행할 수 있는 platform에서는 GPU-side compromise가 host process compromise로 확대될 수 있다.

== Importance ==
이 논문의 중요한 점은 HMM을 단순한 performance/programmability feature가 아니라 새로운 security boundary 변화로 해석한다는 것이다. Prior GPU attack들은 side channel, residual GPU memory disclosure, GPU kernel control-flow hijack을 보여주었지만, host process memory를 직접 corrupt하는 공격은 다루지 못했다. GHOST-ATTACK은 HMM이 이 boundary를 무너뜨릴 수 있음을 보인다.

또한 이 논문은 "address space가 unified되더라도 access authority는 unified되면 안 된다"는 design principle을 제시한다. HMM의 장점은 host pointer를 GPU에서 쓸 수 있게 하는 것이지만, 실제로 GPU kernel이 필요로 하는 것은 모든 host memory가 아니라 kernel argument로 전달된 shared data이다. SHELL은 이 distinction을 compiler와 driver가 함께 복원하는 설계로, future heterogeneous memory system의 security policy를 논의할 때 재사용하기 좋은 framing을 제공한다.

== Main Idea ==
핵심 아이디어는 HMM 환경에서 GPU kernel의 host memory 접근 권한이 host process 전체 address space로 넓어지는 순간, GPU memory-safety bug가 host memory exploit primitive로 승격된다는 것이다. 공격자는 GPU kernel context를 장악한 뒤 host pointer와 libcuda-rt metadata를 이용해 host address layout을 알아내고, return address나 GOT entry 같은 control-relevant host data를 overwrite할 수 있다.

방어의 핵심 아이디어는 GPU가 접근해도 되는 host memory를 "host가 kernel launch 시점에 의도적으로 전달한 shared data"로 한정하는 것이다. SHELL은 Clang/LLVM instrumentation으로 `cudaLaunchKernel`에 전달되는 pointer의 allocation source를 추적해 shared region whitelist를 만들고, NVIDIA GPU driver의 HMM page-fault handler에서 faulting address가 whitelist 안에 있을 때만 page migration을 허용한다.

추가로, NVIDIA HMM page migration이 64KB granularity로 처리되기 때문에 단순 page-fault check만으로는 sub-page leakage가 남는다. SHELL은 shared data와 host-only data가 같은 64KB migration block에 섞이지 않도록 global, heap, stack shared allocation을 별도 aligned region/pool로 옮긴다.

== Design ==
# GHOST-ATTACK threat model
: Victim은 HMM-enabled CUDA application이며, PyTorch/TensorFlow 같은 framework처럼 libcuda-rt를 dynamic link한다. Attacker는 GPU kernel memory-safety bug를 crafted input으로 exploit하거나, WebGPU/HIPscript-like interface를 통해 attacker-supplied GPU kernel을 실행할 수 있다고 가정한다. 목표는 GPU kernel privilege에서 host process memory integrity와 control flow를 compromise하는 것이다.

# ASLR bypass through libcuda-rt
: GHOST-ATTACK은 libcuda-rt가 host process에 남기는 memory layout artifacts를 이용한다. 첫째, GPU launch 시 `nvidiactl` device mapping이 fixed virtual address에 배치되어 reliable anchor가 된다. 둘째, 그 mapping의 fixed offset에 host heap object를 가리키는 pointer가 있다. 셋째, heap에는 libcuda-rt가 관리하는 host address layout information structure가 있어 text, data, heap, stack, shared library mapping base/end를 복원할 수 있다. 이 조합은 GPU kernel이 brute force 없이 host ASLR을 우회하게 만든다.

# Host control-flow hijack
: ASLR 우회 후 GPU kernel은 HMM을 통해 host stack 또는 shared library writable metadata를 직접 쓴다. 논문은 두 variant를 설명한다. 하나는 `cudaDeviceSynchronize()`가 기다리는 동안 libcuda-rt blocking loop의 return address를 stack에서 찾아 gadget address로 바꾸는 방식이다. 다른 하나는 partial RELRO 때문에 writable한 libcuda-rt GOT entry, 예를 들어 `free()` entry를 overwrite하는 방식이다. 이 설계는 GPU memory bug가 host arbitrary code execution으로 이어질 수 있음을 보여준다.

# SHELL static identification of shared data
: SHELL은 GPU kernel이 접근할 legitimate shared data를 GPU kernel 실행 중 access trace에서 추정하지 않는다. 대신 host-side `cudaLaunchKernel()` argument가 shared pointer의 root라는 observation을 사용한다. LLVM/Clang pass가 kernel launch argument에서 backward use-def analysis를 수행하여 global variable, heap allocation, stack allocation, memory mapping 중 GPU에 전달될 수 있는 region을 찾고, 이 region 정보를 GPU driver의 SHELL table에 등록한다. 이 방식은 compromised GPU kernel이 임의 host address를 읽었다는 사실만으로 그 address를 shared data로 오인하는 문제를 피한다.

# Isolated shared allocation
: SHELL은 allocation kind별로 shared data를 host-only data와 분리한다. Global shared data는 64KB-aligned Shared Main Data region으로 옮긴다. Heap shared data는 jemalloc `mallocx` 기반 separate arena에서 할당한다. Stack shared data는 function prologue/epilogue instrumentation을 통해 heap allocation/free로 변환한다. `mmap` 기반 shared allocation은 custom `MAP_HMM` flag를 붙이고 64KB alignment/padding과 driver table update를 수행한다. tradeoff는 compiler/runtime instrumentation과 custom allocator dependency가 생기는 대신, HMM의 sub-page migration granularity 문제를 줄인다는 것이다.

# Runtime access control in GPU driver
: SHELL은 NVIDIA GPU driver의 HMM page-fault path를 enforcement point로 사용한다. GPU가 GPU-only memory를 access하면 기존처럼 GART translation이 성공한다. GPU가 host memory를 access해 page fault가 발생하면, driver는 faulting address가 SHELL table의 shared data whitelist에 있는지 확인한다. Shared data이면 migration을 허용하고, host-only data 또는 unmapped page이면 migration을 거부한다. 이 invariant는 "GPU kernel은 GPU-only data와 explicitly shared host data만 접근한다"는 것이다.

== Result ==
Evaluation setup은 Ubuntu 22.04 LTS, Linux kernel 6.8.0, Intel Core i7-14700K, NVIDIA RTX 4060, open-source NVIDIA driver 570, CUDA 12.8이다. 공격 실험에서는 Python 3.12 process와 Chrome 138.0.7191 WebGPU platform을 사용했다.

공격 결과는 두 scenario로 제시된다. 첫째, PyTorch-like vulnerable GPU kernel에 fixed-size device buffer overflow를 넣고 attacker payload SASS를 input에 포함시켰다. GPU-side control flow가 attacker payload로 넘어간 뒤, payload는 host address layout을 leak하고 libcuda-rt의 `free()` GOT entry를 `system()`을 호출하는 gadget으로 바꾸어 arbitrary command execution을 달성했다. 둘째, Chrome WebGPU scenario에서는 현재 WebGPU가 CUDA general-purpose kernel을 직접 지원하지 않기 때문에, 저자들이 malicious GPU kernel을 Chrome WebGPU platform 내부 rendering function에 직접 embedding했다. 이 kernel은 Chrome GPU process의 host address layout과 libcuda-rt GOT entry를 leak하고, `ioctl` gadget을 호출하도록 GOT를 바꾸었다.

방어 결과에서 SHELL은 vulnerable GPU kernel 자체의 exploit이나 GPU-side payload execution을 막지는 못하지만, payload가 host-only data를 접근하는 것을 막는다. 그 결과 host address layout leak과 host GOT overwrite가 차단된다. Chrome/WebGPU case에서는 future GPU kernel launching module을 가정하고 그 module에 SHELL instrumentation을 적용했을 때 attacker-controlled GPU kernel의 host-only memory access가 차단되었다고 보고한다.

Performance evaluation은 NVIDIA의 HMM-enabled ERA5 climate dataset processor sample을 대상으로 했다. Baseline jemalloc version의 elapsed time은 3.1088초이고 SHELL-enabled version은 3.1398초로, end-to-end overhead는 0.9%이다. `nvidia-smi` 기준 GPU memory allocation size는 baseline과 SHELL 모두 5,824 MiB로 같아, 64KB migration block alignment가 이 workload에서 extra fragmentation을 만들지 않았다고 보고한다.

Microbenchmark에서는 shared data allocation 중 `mmap`이 3회(0.2%), `malloc`이 1,440회(99.8%)였다. Custom `mmap + ioctl` latency는 5.512us에서 6.106us로 10.7% 증가했지만 호출 수가 작아 end-to-end 영향은 작다. GPU driver page-fault handler latency는 97.02us에서 108.75us로 12.1% 증가했다. 이 결과는 SHELL overhead가 enforcement point에서는 관측되지만, HMM workload 전체 실행 시간에는 작게 반영된다는 주장을 뒷받침한다.

== Contribution ==
# HMM-enabled CUDA system에서 GPU kernel이 host memory integrity를 직접 compromise할 수 있는 GPU-to-host attack surface를 정식 threat model로 제시했다.
# libcuda-rt의 fixed mapping, heap pointer exposure, host layout metadata를 이용해 GPU kernel에서 host ASLR을 우회하는 concrete GHOST-ATTACK primitive를 보였다.
# PyTorch-like vulnerable GPU kernel과 Chrome WebGPU prototype scenario에서 GPU-originated attack이 host control-flow hijack으로 이어질 수 있음을 실험적으로 보였다.
# LLVM/Clang static instrumentation과 NVIDIA GPU driver page-fault enforcement를 결합한 SHELL을 제안해 shared data와 host-only data 사이의 access-control boundary를 복원했다.
# ERA5 HMM sample application과 microbenchmark를 통해 SHELL이 공격 variant를 차단하면서 end-to-end overhead 0.9% 수준으로 동작함을 보였다.

== Criticisms ==
HMM 자체가 아직 production application에서 널리 쓰이지 않기 때문에, performance evaluation의 real-world scope는 제한적이다. 논문도 이를 인정하고 NVIDIA HMM sample인 ERA5 processor를 대표 workload로 사용한다. 더 다양한 ML framework, multi-process service, long-running inference workload에서 allocation pattern과 page-fault behavior가 어떻게 달라지는지는 추가 검증이 필요하다.

Compiler instrumentation은 host code가 분석 가능하고, shared data가 `cudaLaunchKernel` argument path를 통해 식별된다는 가정을 둔다. Dynamic loading, closed-source library, inline assembly, driver-level custom launch path, unusual pointer aliasing이 있을 때 false negative 또는 compatibility issue가 생길 수 있는지 evaluation이 충분히 크지는 않다.

SHELL은 GPU kernel의 memory-safety vulnerability 자체를 제거하지 않는다. 공격 payload가 host-only memory에 접근하는 것을 막는 boundary defense이므로, GPU memory 내부 data corruption, GPU-side control-flow hijack, denial-of-service는 별도 방어가 필요하다.

== Conclusion ==
이 연구는 HMM을 단순한 unified memory convenience가 아니라 host/GPU privilege boundary를 바꾸는 security-critical abstraction으로 바라보게 만든다. GHOST-ATTACK은 GPU kernel memory bug나 attacker-supplied GPU code가 host ASLR bypass와 host control-flow hijack으로 이어질 수 있음을 보였고, SHELL은 shared data identification과 driver page-fault enforcement를 결합해 HMM의 programmability를 유지하면서 host-only memory access를 막을 수 있음을 보였다. 따라서 heterogeneous memory management, GPU security, compiler-assisted runtime isolation을 연결하는 신선한 관점을 제시한다.

[[분류: IEEE S&P]]
[[분류: GPU]]

Forest: Access-aware GPU UVM Management

2026-06-19T02:01:50Z

Noribot: Nori: update draft.md

{{Paper
|title=Forest: Access-aware GPU UVM Management
|author=Mao Lin, Yuan Feng, Guilherme Cox, Hyeran Jeon
|conference=ISCA
|year=2025
}}

== 개요 ==
이 논문은 [[GPU Unified Virtual Memory]]에서 기존 [[TBNp|Tree-based Neighboring Prefetcher]]가 access pattern을 모른 채 동일한 tree 설정을 모든 메모리 블록에 적용할 때 왜 불필요한 page migration과 page thrashing이 발생하며, data object별 access pattern에 맞춘 heterogeneous TBNp configuration으로 이를 어떻게 줄일 수 있는지를 다룬다.
GPU 하드웨어의 page access counter를 access recency 추적용으로 재해석하고, UVM driver가 그 정보를 이용해 prefetch tree와 eviction policy를 동적으로 조정한다.

== Motivation ==
[[Unified Virtual Memory]]는 CPU memory를 GPU memory의 확장 공간처럼 사용할 수 있게 해 주지만, GPU가 CPU memory에 있는 page를 처음 접근할 때 far-fault handling과 page migration이 critical path에 들어간다. 특히 modern GPU workload와 [[Deep Learning]] model은 memory footprint가 커지고 있어 memory oversubscription 상황에서 UVM 성능이 중요해진다.

; On Configuration Doesn't Fit All
: NVIDIA GPU 계열에서 쓰이는 TBNp는 2MB VABlock을 full binary tree로 관리하고 64KB leaf node 단위로 migration한다. 이 방식은 locality가 강한 workload에서는 far-fault를 줄일 수 있지만, 모든 application과 모든 data object에 동일한 tree size와 leaf size를 적용한다. 논문은 실험적으로 15개 workload 중 baseline configuration인 2MB tree, 64KB leaf가 최적인 경우가 없다고 보고한다. 또한 같은 application 안에서도 kernel과 data object마다 LS, HCHI, HCLI, LC 같은 access pattern이 달라 하나의 homogeneous prefetcher로는 충분하지 않다.

; TBNp, A Hidden Source of Memory Oversubscription
: 또한 paper는 [[Prefetching]] accuracy와 timeliness뿐 아니라 GPU memory oversubscription에서 prefetcher가 memory pressure를 악화시킬 수 있음을 보여준다. Figure 6의 분석에 따르면 일부 application에서는 migrated page의 5%에서 48%가 eviction 전까지 전혀 access되지 않았고, 일부 workload에서는 반복 thrashing page 수가 memory footprint의 5.7x를 넘는다. 이 framing은 future UVM work에서 prefetcher와 eviction policy를 분리해 볼 수 없다는 근거로 유용하다.

; Driver-driven Page Eviction and Trashing
: 또 다른 문제는 eviction이다. 기존 UVM driver의 LRU는 실제 GPU-side access recency가 아니라 far-fault event 순서에 기반한다. 따라서 device에서 최근에 많이 쓰인 page라도 fault history 기준으로 오래된 것으로 보이면 eviction될 수 있고, 이후 다시 migration되며 page thrashing이 발생한다.

=== [[TBNp]] ===
본 논문은 TBNp를 버리지 않고, TBNp의 tree 구조를 access-aware하게 바꾸는 방향을 택한다. Forest의 문제의식은 다음과 같다.

# TBNp 자체는 UVM far-fault를 줄이는 데 효과적이다.
# 하지만 fixed 2 MB tree와 64 KB leaf node는 모든 access pattern에 맞지 않는다.
# UVM driver만으로는 GPU-side access pattern을 충분히 볼 수 없기 때문에, TBNp가 불필요한 migration과 thrashing을 만들 수 있다.
# 따라서 data object별 access pattern을 감지하고, TBNp tree size와 leaf size를 object별로 다르게 설정해야 한다.

== Main Idea ==
핵심 아이디어는 UVM-managed data object마다 access pattern이 다르므로, 고정된 TBNp tree 하나를 쓰지 말고 object별로 tree size와 leaf size를 바꾸자는 것이다. Forest는 GPU-side access counter를 이용해 page access 순서를 관찰하고, UVM driver가 이를 네 가지 access pattern으로 분류한 뒤 각 pattern에 맞는 tree configuration을 선택한다.

Forest에서 recency가 중요한 이유는 기존 UVM driver가 실제 GPU access recency를 잘 모르기 떄문이다. 기존 driver는 page fault 순서를 기준으로 LRU 비슷하게 판단하는데, page fault가 안 난 page라도
GPU에서 계속 hot하게 접근 중일 수 있다. Forest의 ATT는 page access counter를 이용해서 “이 page/object가 실제로 언제 접근됐는가”를 기록하게 하였다. 즉 access frequency가 “몇 번 접근됐나”라면, access recency는 “마지막으로 언제
접근됐나”에 가깝다. 이를 통해서 '''Hardware-level의 LRU를 최적으로 구할 수 있다.'''

ATT를 이용해서 얻은 정보로, Pattern과 configuration의 대응을 다음과 같이 할 수 있다. 기존에는 분류를 Regularity로만 판단했는데, 본 논문에서는 Intensity도 축으로 두어, 더 Fine-grained한 분류를 하였다.

# Linear/Streaming (LS): page가 순차적으로 접근되고 재사용이 낮다. 큰 tree와 큰 leaf를 사용해 aggressive prefetching을 한다.
# Non-Linear High-Coverage High-Intensity (HCHI): 넓은 address range에서 많은 page가 빠르게 접근되지만 linear하지 않다. 작은 tree와 기본 leaf로 prefetch 범위를 제한한다.
# Non-Linear High-Coverage Low-Intensity (HCLI): 넓은 range를 sparse하게 접근한다. 작은 tree와 작은 leaf를 사용해 불필요한 migration을 줄인다.
# Non-Linear Low Coverage (LC): coverage가 낮거나 명확한 pattern이 없다. default TBNp configuration을 사용한다.

SpecForest는 Forest의 profiling delay를 줄이기 위한 확장이다. 반복 kernel의 이전 pattern을 재사용하고, compiler static analysis로 fixed-stride LS access를 미리 표시하며, indirect indexing이 비슷한 data object들을 similarity group으로 묶어 한 object의 classification 결과를 같은 group에 전파한다.

== Design ==
# Access Time Tracker (ATT): GPU GMMU 쪽에 object table을 두고 UVM object별 VPN range, access timer, recency order, cease bit를 관리한다. 기존 page access counter가 access frequency를 담는 대신, Forest에서는 object-local access timer 값을 page counter에 기록해 page access order를 나타내도록 한다. 이 설계는 device-side access recency를 driver가 읽을 수 있게 만드는 핵심 하드웨어 변경이다.
# Access Pattern Detector (APD): UVM driver module로, profiling interval마다 access counter 정보를 가져와 object별 pattern을 분류한다. LS는 page number와 access time의 linear regression에서 R^2 threshold를 이용해 판정하고, HCHI/HCLI/LC는 accessed VPN range coverage와 accessed page count intensity threshold로 구분한다. pattern이 결정되면 object의 profiling을 멈추고 pattern table에 기록한다.
# Prefetch Engine (PE) 확장: 기존 UVM driver의 TBNp traversal과 migration path를 유지하되, non-leaf node마다 isolation bit와 motion bit를 추가한다. isolation bit는 child subtree를 서로 독립된 prefetch tree로 분리해 tree size를 조절하고, motion bit는 child nodes를 하나의 basic block처럼 취급해 leaf size를 조절한다. 이 두 bit로 16KB 단위의 기본 partition 위에서 object별 heterogeneous tree를 만든다.
# Access-aware eviction: 기존 far-fault 기반 LRU 대신 ATT의 object recency order와 page access counter의 recency 값을 이용한다. 먼저 가장 오래된 object를 찾고, 그 object 안에서 가장 오래된 page가 포함된 leaf node를 evict한다. global memory 전체를 훑는 대신 object 단위로 search space를 줄이면서 GPU-side recency를 반영한다.
# SpecForest: pattern recording, static LS detection, access similarity detection을 추가한다. compiler가 UVM object의 simple stride access 또는 동일 index expression 기반 similarity group을 표시하고, driver는 이 정보를 cudaMallocManaged flag 확장 형태로 받아 초기 tree configuration을 앞당긴다.

=== Bringing It All Together ===
# Kernel launch 시점에 UVM driver가 해당 kernel이 접근할 managed memory object들의 VPN range를 ATT object table에 등록한다.
# Kernel이 실행되는 동안 GPU memory access는 GMMU에 붙은 ATT의 access timer와 recency order를 갱신한다.
# 어떤 object의 access count가 profiling threshold에 도달하면 ATT가 driver에 interrupt를 보내고, driver는 기존 access counter copy path를 이용해 object의 page access timing 정보를 가져온다.
# APD가 이 정보를 LS/HCHI/HCLI/LC 중 하나로 분류한다. 만약 여러 번 profiling해도 pattern이 결정되지 않으면 Forest는 해당 object를 default LC pattern으로 처리한다.
# Pattern이 결정되면 driver는 pattern table에 결과를 기록하고, PE는 해당 object의 TBNp tree에 motion bit와 isolation bit를 설정해 tree size와 leaf size를 바꾼다.
# Driver는 ATT의 cease bit를 set하여 이미 pattern이 결정된 object에 대해서는 추가 profiling interrupt를 멈춘다.
# 이후 page fault가 발생하면 UVM driver는 새로 configured된 tree를 사용해 fault handling과 prefetching을 수행한다.
# Oversubscription 때문에 eviction이 필요하면 ATT object table에서 가장 오래된 object를 찾고, 그 object 내부에서 가장 오래된 page가 포함된 leaf node를 evict한다.
# 즉 Forest의 runtime loop는 '''ATT가 access recency를 기록하고, APD가 pattern을 결정하며, PE가 object별 TBNp를 재구성하고, eviction policy가 같은 recency 정보를 재사용하는 구조'''이다.

== Result ==
평가는 UVM-enabled GPGPU-Sim 4.0과 UVM emulator를 확장해 수행했다. 기본 설정은 NVIDIA Turing-style GPU, 4KB page, 45us far-fault handling latency, PCIe 3.0 x16, 150% memory oversubscription이다. benchmark는 UVMBench, InterplayUVM, Tango 기반 15개 workload이며, 추가로 Accel-Sim 통합 환경에서 AlexNet, ResNet50, BERT, Whisper를 평가했다.

주요 결과는 다음과 같다.

# Overall speedup: Forest는 baseline TBNp 대비 평균 1.72x, SpecForest는 1.86x speedup을 보인다. 논문 초록과 introduction은 SpecForest가 state-of-the-art 대비 최대 1.39x 수준의 향상을 보인다고 요약한다.
# Far-fault reduction: linear workload에서는 큰 tree가 2MB boundary를 넘는 sequential access의 far-fault를 줄인다. mixed-pattern workload에서는 HCHI/HCLI object에 작은 tree와 leaf를 적용해 불필요한 migration과 thrashing을 줄인다.
# Thrashing reduction: optimal tree configuration만 적용해도 page thrashing이 평균 25% 줄고, access-aware LRU를 추가하면 추가로 7% 감소한다.
# Profiling overhead reduction: SpecForest는 평균 profiling step을 223회에서 10회로 줄인다. pattern recording이 특히 mixed-pattern benchmark에서 효과가 크고, static analysis와 similarity detection은 각각 평균 2% 수준의 추가 speedup을 준다.
# Sensitivity: oversubscription 125%에서 200%까지 SpecForest는 baseline 대비 1.57x에서 1.95x speedup을 유지한다. Pascal, Volta, Turing, Ampere, Hopper-style GPU configuration에서도 이득이 유지된다.
# Real-world DL: AlexNet, ResNet50, BERT, Whisper에서 SpecForest는 평균 1.51x, 최대 1.62x speedup을 보인다. CNN은 LS가 지배적이고, Transformer 계열은 self-attention의 irregular dense access 때문에 HCHI 비중이 커져 per-object heterogeneous prefetching의 가치가 커진다.

이 결과는 성능 개선이 단순한 prefetch aggressiveness 증가가 아니라, access pattern에 맞춘 tree shape와 GPU-side recency 기반 eviction이 far-fault와 memory thrashing을 동시에 줄인 결과라는 논문의 thesis를 뒷받침한다.

== Contribution ==
# TBNp의 homogeneous configuration이 unnecessary migration과 page thrashing을 유발한다는 문제를 workload, kernel, data object 수준에서 분석했다.
# GPU page access counter를 access recency tracker로 재해석하는 ATT와 UVM driver 기반 APD를 결합해 object-level access pattern detection을 제안했다.
# isolation bit와 motion bit로 기존 TBNp semantics를 유지하면서 tree size와 leaf size를 object별로 조정하는 heterogeneous TBNp mechanism을 설계했다.
# access counter 기반 pseudo-LRU eviction으로 far-fault history가 아닌 실제 GPU-side access recency를 eviction에 반영했다.
# SpecForest를 통해 compiler-assisted static detection, pattern recording, similarity group propagation으로 runtime profiling delay를 줄였다.
# general-purpose GPU benchmark와 real-world DL workload에서 Forest/SpecForest가 baseline 및 prior UVM optimization보다 높은 성능을 보임을 시뮬레이션으로 평가했다.

== Criticisms ==
# 하드웨어 수정이 필요하기 때문에, 현재 Production-level에서 사용되는 Set-up에서는 사용할 수 없다.
# 조금 너무 어렵게 논문이 쓰인 느낌이 있다. Top-down방식으로 Componenet-by-componenet으로 서술하면 좀더 쉽게 설명할 수 있을 것 같다는 생각이 든다. 예를 들어서 굳이 APD나 ATT혹은 Prefetcher Engine은 너무 Classficiation하기 위한 용어처럼 느껴진다 (E.g., Speculative Forest -> Compiler Optimization 으로 변경, APD -> Forest APIs 로 변경, Overview 논문 Design 섹션전에 삽입.. 등). 또한 중복되는 설명이 많다.
# Object의 access pattern이 runtime 중간에 dynamic하게 바뀌는 경우에는 둔감할 수 있다. Forest는 APD가 object의 pattern을 한 번 결정하면 ATT의 cease bit를 set하여 해당 object에 대한 추가 profiling interrupt를 멈춘다. 따라서 같은 kernel execution 안에서 object가 처음에는 LS처럼 동작하다가 이후 sparse/irregular pattern으로 바뀌는 intra-kernel phase change는 잘 따라가지 못할 수 있다. 이는 profiling overhead를 줄이기 위한 early classification의 tradeoff이다.

== [[Conclusion]] ==
Forest는 GPU UVM의 병목을 "page fault를 얼마나 빨리 처리하는가"보다 "어떤 data object를 어떤 단위로 미리 가져오고 언제 내보내는가"의 문제로 본다. 이 관점에서 TBNp의 tree 구조를 고정 정책이 아니라 access pattern별 policy substrate로 바꾸고, device-side recency를 driver decision에 연결한다.

나중에 기억할 점은 Forest가 prefetcher를 더 aggressive하게 만든 논문이 아니라는 것이다. LS에는 aggressiveness를 키우고, irregular sparse/high-coverage access에는 aggressiveness를 줄이며, eviction도 실제 access recency에 맞춘다. 즉, UVM oversubscription에서 prefetching과 eviction은 같은 memory pressure budget을 공유한다는 점을 설계로 보여준 paper이다.

== Assumptions and Verification ==

* Metadata는 PDF extracted text의 ACM reference format과 첫 페이지 정보를 기준으로 작성했다.
* PDF source는 `/home/jhyohan/MPDK/mBPF-Usecase/.nori/runs/focused-signal/sources/3695053.3731047.pdf`, extracted text는 같은 directory의 `3695053.3731047.txt`를 사용했다.
* Figure/table의 일부 텍스트는 `pdftotext` extraction에서 layout noise가 있었으므로, 수치와 mechanism은 본문 문장과 caption 중심으로 확인했다.
* MediaWiki 검색 connector는 제공되지 않았으므로 link 이름은 noriwiki에 있을 가능성이 높은 reusable concept 이름으로 작성했다.

[[분류: ACM ISCA]]
[[분류: GPU UVM]]

Memory Harvesting in Multi-GPU Systems with Hierarchical Unified Virtual Memory

2026-06-18T04:59:51Z

Noribot: Nori: update draft.md

{{Paper
|title=Memory Harvesting in Multi-GPU Systems with Hierarchical Unified Virtual Memory
|author=Sangjin Choi, Taeksoo Kim, Jinwoo Jeong, Rachata Ausavarungnirun, Myeongjae Jeon, Youngjin Kwon, Jeongseob Ahn
|conference=2022 USENIX Annual Technical Conference (USENIX ATC 22)
|year=2022
}}

== 개요 ==
이 논문은 shared multi-GPU server에서 어떤 [[GPU]]는 memory oversubscription으로 host memory를 느리게 쓰는 반면, 다른 GPU에는 작은 idle memory가 남는 imbalance가 왜 발생하며, neighbor GPU의 spare memory를 [[GPU Unified Virtual Memory]] 계층에 넣어 이를 어떻게 완화할 수 있는지를 다룬다.

== Motivation ==
GPU 수요가 커지면서 연구실과 산업 환경에서는 multi-GPU server를 여러 job이 공유하는 방식이 일반적이다. DNN training은 batch size를 GPU memory에 거의 맞추어 조정하고, graph analytics는 graph size에 따라 memory demand가 크게 달라진다. 이 때문에 한 서버 안에서도 어떤 GPU는 oversubscription으로 host DRAM을 swap-like backing store로 쓰는 반면, 다른 GPU에는 수백 MB에서 수 GB의 idle memory가 남을 수 있다.

기존 [[GPU Unified Virtual Memory]]는 GPU memory보다 큰 working set을 host memory까지 활용해 실행하게 해 주지만, host memory 접근은 PCIe path를 거친다. 논문은 AWS p3.8xlarge의 NVIDIA V100 환경에서 2 MB migration을 측정했을 때 PCIe는 12.3 GB/s, 16.7 us인 반면 NVLink는 40.1 GB/s, 5.1 us로 약 3배 빠르다고 보고한다. 즉 같은 서버 안의 neighbor GPU memory는 host memory보다 작고 동적으로 변하지만, 접근 latency와 bandwidth 면에서는 훨씬 나은 중간 계층이 될 수 있다.

따라서 문제는 단순히 "GPU memory가 부족하다"가 아니라 "multi-GPU server 전체에는 잠시 놀고 있는 GPU memory가 있는데, 기존 UVM은 이를 capacity extension으로 쓰지 못한다"는 점이다. 이 논문은 이 memory fragmentation을 driver-level memory management 문제로 재구성한다.

== Importance ==
이 논문이 중요한 이유는 GPU memory oversubscription을 single-GPU local memory와 host memory 사이의 문제로만 보지 않고, shared multi-GPU server 전체의 memory hierarchy 문제로 확장했기 때문이다. 이전 연구들은 주로 host memory prefetch/pre-eviction, framework-guided tensor swapping, hardware/runtime co-design, memory compression에 초점을 두었다. 반면 이 논문은 commodity multi-GPU server에 이미 존재하는 NVLink/NVSwitch path와 neighbor GPU idle memory를 사용한다.

또한 memHarvester는 spare memory가 충분하지 않아도 유용하다는 점을 보인다. neighbor GPU memory를 모든 evicted data의 최종 거주지로 쓰는 것이 아니라, host access latency를 critical path에서 숨기는 victim cache이자 staging buffer로 사용한다. 이 framing은 이후 [[ARIADNE: Adaptive UVM Management for Efficient GPU Memory Oversubscription]] 같은 UVM oversubscription 연구와도 보완적이다. ARIADNE가 VABlock locality와 Zero-copy placement에 집중한다면, 이 논문은 multi-GPU topology와 path diversity를 활용하는 방향을 제시한다.

== Background ==
; GPU UVM
: [[GPU Unified Virtual Memory]]는 GPU와 CPU memory를 하나의 virtual address space로 보이게 하고, page fault를 통해 필요한 page를 GPU memory로 migrate한다. GPU memory가 부족하면 기존 GPU-resident page를 host memory로 evict한 뒤 faulted page를 가져와야 하므로, eviction과 fetch가 page fault critical path에 들어간다.

; 2 MB chunk
: 논문 기준 NVIDIA UVM은 GPU physical memory를 2 MB chunk 단위로 관리한다. Page fault 자체는 host architecture의 base page granularity와 연결되지만, eviction은 2 MB chunk 단위로 수행된다. memHarvester는 이 granularity를 유지하면서 neighbor GPU spare chunk를 harvested memory로 표시한다.

; Hierarchical Unified Virtual Memory
: HUVM은 local GPU memory, neighbor GPU spare memory, host memory로 구성된 hierarchy이다. Local GPU에 free space가 없을 때 바로 host memory로 evict하는 대신, NVLink로 연결된 neighbor GPU spare memory를 먼저 victim buffer로 사용한다.

== Challenge ==
# Effective harvesting: neighbor GPU의 idle memory는 작고, spotty하며, workload phase에 따라 변한다. 따라서 spare memory가 전체 overcommitted data를 담지 못해도 성능 이득을 내야 한다.
# Minimal interference: harvested memory를 제공하는 yielding GPU도 application을 실행 중이다. Borrowed memory capacity뿐 아니라 NVLink, PCIe, memory bandwidth 사용이 yielding workload를 방해할 수 있다.
# Low overhead: UVM은 page fault handling과 page table update 비용을 이미 갖고 있다. HUVM이 추가 metadata와 background work를 넣으면 NVLink/PCIe bandwidth를 충분히 활용하기 전에 software overhead가 병목이 될 수 있다.
# Framework-agnostic: shared server의 jobs는 PyTorch DNN training, cuGraph graph analytics처럼 서로 다른 framework를 쓴다. 따라서 solution은 application/framework modification 없이 driver layer에서 동작해야 한다.

== Main Idea ==
=== Characterization ===
논문의 관찰은 shared multi-GPU server에서 memory pressure가 GPU별로 균등하지 않다는 것이다. Figure 1의 consolidation scenarios에서는 PageRank, BFS, WCC, Louvain 같은 graph workload가 host memory를 쓰는 동안 VGG16, MobileNet, ResNet101 또는 partitioned PageRank가 실행되는 다른 GPU에는 idle memory가 남는다. 기존 UVM은 이 idle memory를 보지 못하므로 host memory를 느린 backing store로 사용한다.

=== Optimization ===
핵심 아이디어는 neighbor GPU의 spare memory를 local GPU와 host memory 사이의 fast victim cache로 추가하는 것이다. Spare memory가 모든 evicted page를 보관할 만큼 크지 않아도, 최근 evicted chunk를 NVLink-attached memory에 잠시 두고 background writeback과 prefetch를 병렬화하면 host memory access latency를 critical path 밖으로 밀어낼 수 있다.

memHarvester는 이 아이디어를 centralized driver-level coordinator로 구현한다. It dynamically harvests spare 2 MB chunks, evicts local GPU chunks to harvested memory, writes them back to host in the background, and reclaims removable chunks quickly when the yielding GPU needs memory again. 즉 harvested memory는 빌린 capacity이면서도 즉시 돌려줄 수 있어야 하는 transient resource이다.

== Design ==
# HUVM data path
#: Local problem: stock UVM은 local GPU memory가 부족하면 host memory로 evict하고, 이후 재접근 시 PCIe를 통해 fetch한다.
#: Mechanism: HUVM은 local GPU, harvested neighbor GPU memory, host memory를 hierarchy로 구성한다. Evicted 2 MB chunks can reside in harvested memory first, and later be copied back to host.
#: Why it helps: NVLink path가 PCIe host path보다 빠르므로 eviction latency와 refetch latency를 줄일 수 있다.
#: Tradeoff: harvested memory는 yielding GPU의 resource이므로 capacity와 bandwidth가 동적으로 변하고, 다른 harvester와 공유될 수 있다.

# Spare memory management and reclamation
#: Local problem: yielding GPU가 다시 memory를 필요로 하면 borrowed chunks를 빠르게 반환해야 한다.
#: Mechanism: memHarvester는 GPU별 free 2 MB chunk list를 보고 spare memory를 식별한다. Harvested chunk metadata를 표시하고, evicted list와 removable list를 관리한다. Evicted chunk가 host에 background writeback되면 removable로 표시되어 yielding GPU가 fault를 낼 때 먼저 회수된다.
#: Why it helps: harvested memory를 "언젠가 돌려줄 cache"가 아니라 "host copy가 준비되면 즉시 회수 가능한 cache"로 만들어 interference를 줄인다.
#: Tradeoff: writeback이 완료되기 전에는 해당 chunk를 즉시 회수할 수 없으므로, writeback throughput과 large-page support가 중요하다.

# Eviction path optimization
#: Local problem: demand fault가 발생한 뒤 eviction을 시작하면 eviction time이 page fault critical path에 들어간다.
#: Mechanism: free chunk가 threshold 아래로 내려가면 memHarvester가 pre-eviction thread를 실행한다. Victim은 stock UVM과 같은 LRU 계열 policy로 고르고, target yielding GPU는 round-robin으로 선택한다. Host writeback에는 2 MB large page를 사용해 512개의 4 KB page population을 하나의 2 MB operation으로 줄인다.
#: Why it helps: pre-eviction은 future demand fault 전에 free chunk를 확보하고, large page eviction은 host-side writeback과 removable chunk 생성 속도를 높인다.
#: Tradeoff: pre-eviction은 잘못된 victim을 미리 내보낼 수 있고, round-robin target selection은 topology-aware optimal policy라기보다 hotspot avoidance에 가까운 heuristic이다.

# Parallel fault handling
#: Local problem: UVM은 fault batch를 처리하지만, host memory에서 faulted pages를 순차적으로 가져오면 PCIe latency가 남는다.
#: Mechanism: memHarvester는 fault batch의 head는 local GPU로, tail은 harvested memory로 처리하는 별도 kernel thread를 둔다. 서로 다른 GPU의 PCIe lane을 병렬로 활용하고, harvested memory에 들어간 page는 이후 NVLink로 local GPU에 들어온다.
#: Why it helps: batch 안의 page fetch를 local path와 yielding GPU path로 나누어 host fetch latency 일부를 숨긴다.
#: Tradeoff: fault batch coordination에 mutex와 metadata update가 필요하며, harvested memory에 미리 올린 page가 실제로 곧 쓰이지 않으면 이득이 줄어든다.

# Multi-path parallel prefetcher
#: Local problem: spare memory가 일부 evicted chunks만 담고 있을 때, host memory access를 완전히 피할 수는 없다.
#: Mechanism: memHarvester는 page fault history에서 next-line/stride pattern을 추출하고 기본 32 MB를 prefetch한다. Harvested memory에 있는 chunk는 NVLink로 local GPU에 prefetch하고, host memory에 있는 chunk는 active harvester 수와 PCIe congestion에 따라 harvested memory 또는 local GPU로 prefetch한다.
#: Why it helps: NVLink와 PCIe path를 동시에 사용하고, host data를 neighbor GPU spare memory에 미리 올려 다음 fault의 source를 host에서 harvested memory로 바꾼다.
#: Tradeoff: 여러 harvester가 하나의 yielding GPU를 공유하면 yielding GPU의 PCIe lane이 bottleneck이 될 수 있다. 이 경우 memHarvester는 host-to-local prefetch를 선택해 contention을 피하지만, harvested-memory prefetch 이득은 줄어든다.

# Driver-transparent implementation
#: Local problem: shared GPU server의 workload는 framework와 application이 다양하므로 user-level rewrite가 어렵다.
#: Mechanism: prototype은 NVIDIA UVM driver 460.67 위에 1,838 SLOC의 C code로 구현되었다. PyTorch와 cuGraph workload를 application modification 없이 실행한다.
#: Why it helps: UVM driver가 GPU별 page table, residency metadata, fault stream을 볼 수 있으므로 multi-GPU memory movement를 centralized하게 조정할 수 있다.
#: Tradeoff: NVIDIA UVM driver 내부 수정이 필요하므로 production deployment는 driver version, kernel memory allocation, GPU topology 지원에 의존한다.

== Result ==
평가는 AWS p3.8xlarge에서 수행되었다. 이 서버는 16 GB NVIDIA V100 GPU 4개, NVSwitch/NVLink 2.0 interconnect, PCIe 3.0으로 연결된 240 GB host memory를 갖는다. Baseline은 stock UVM인 Base와 host memory에 대한 pre-eviction/prefetch prior approach를 모사한 Pre-ef-host이다. Workload는 PyTorch 1.10.1 DNN training과 cuGraph 21.12 graph analytics를 사용한다.

Inter-job harvesting에서는 graph workload가 다른 GPU의 DNN 또는 graph workload가 남긴 spare memory를 사용한다. Case-1에서 PageRank는 VGG16과 WCC가 남긴 총 4.64 GB spare memory를 harvest한다. 이 spare memory는 PageRank의 overcommitted memory 13.92 GB보다 작지만, memHarvester는 Base 대비 3.53배, Pre-ef-host 대비 1.31배 speedup을 보인다. 이는 harvested memory가 전체 overflow를 담지 못해도 host access latency를 숨길 수 있음을 보여준다.

Case-2에서 BFS는 MobileNet과 ResNet101이 실행되는 GPU의 spare memory를 사용해 Base 대비 3.52배, Pre-ef-host 대비 2.1배 빨라진다. 이때 yielding workload의 performance impact는 약 7-9%이다. Case-3에서는 WCC와 BFS가 PageRank가 실행되는 두 GPU의 spare memory를 공유하며, WCC는 Base 대비 3.83배와 Pre-ef-host 대비 2.71배, BFS는 Base 대비 3.21배와 Pre-ef-host 대비 2.67배 speedup을 보인다. 논문이 주장하는 "prior approach 대비 up to 2.71x"는 이 case에서 나온다. Case-4에서는 WCC와 Louvain이 ResNet101이 남긴 4.16 GB spare memory 하나를 공유하며, Base 대비 약 2.1-2.36배, Pre-ef-host 대비 30-40% 개선을 보인다.

Ablation은 각 mechanism이 다른 병목을 줄인다는 점을 뒷받침한다. Harvesting(H)은 NVLink victim cache로 eviction/fetch latency를 줄이고, pre-eviction(PE)은 eviction을 demand fault critical path에서 빼며, large page(LP)는 host writeback을 빠르게 만들어 removable chunk 생성을 촉진한다. Parallel fetch(PLF)는 fault batch 처리 latency를 줄이고, multi-path parallel prefetcher(MPF)는 local-only prefetch보다 큰 이득을 보인다. 다만 Case-4처럼 하나의 yielding GPU에 여러 harvester가 몰리면 MPF는 yielding GPU PCIe contention을 피하기 위해 host-to-local prefetch를 선택하므로 local prefetcher 대비 추가 이득이 작다.

Sensitivity study에서는 2 MB stride와 32 MB prefetch size를 기본값으로 선택한다. Available spare memory를 5%에서 60%까지 바꾸는 실험에서는 5%, 즉 800 MB의 spare memory만 있어도 네 개 graph workload 모두에서 2배 이상 성능 향상이 나온다. 어느 정도 spare memory가 active working set을 수용하면 추가 spare memory를 늘려도 성능은 더 좋아지지 않는다.

Interference 측정에서는 ResNet101이 yielding GPU에서 실행되고 VectorAdd microbenchmark가 harvested memory traffic을 만든다. Harvester가 하나일 때 ResNet101 degradation은 최대 3%로 작지만, 세 harvester가 동시에 접근하면 degradation이 13%까지 커진다. 논문은 이는 memory bandwidth와 NVLink traffic contention 때문이며, throttling은 future work로 남긴다.

Intra-job harvesting에서는 pipeline parallel DNN training처럼 하나의 multi-GPU job 내부에서 GPU별 memory demand가 불균형한 경우를 다룬다. GNMT16과 ResNet50에서는 memHarvester throughput이 Pre-ef-host보다 1.5-1.6배 높다. GNMT8과 VGG16처럼 yielding GPU들의 aggregate spare memory가 충분하지 않은 경우에도 Pre-ef-host 대비 각각 2.16배, 1.24배 throughput improvement를 보인다.

== Contribution ==
# Shared multi-GPU server에서 GPU별 memory pressure와 idle memory가 동시에 발생하는 memory imbalance를 DNN training과 graph analytics consolidation scenario로 정량화하였다.
# Local GPU memory와 host memory 사이에 neighbor GPU spare memory를 넣는 Hierarchical Unified Virtual Memory(HUVM) abstraction을 제안하였다.
# Dynamic spare memory harvesting, background writeback, fast reclamation을 결합해 yielding GPU의 memory를 transient victim cache로 사용하는 memHarvester를 설계하였다.
# Pre-eviction, 2 MB large page eviction, parallel fault-batch fetch, multi-path parallel prefetcher를 통해 eviction latency와 fetch latency를 각각 줄였다.
# NVIDIA UVM driver 수정만으로 PyTorch와 cuGraph workload에 transparent하게 적용 가능한 prototype을 구현하였다.
# Inter-job과 intra-job harvesting 실험을 통해 Base 및 host-only pre-eviction/prefetch baseline 대비 성능 향상과 yielding workload interference를 함께 제시하였다.

== Conclusion ==
이 연구는 GPU memory oversubscription을 단일 GPU의 capacity 부족 문제가 아니라 shared multi-GPU server의 memory fragmentation과 hierarchy design 문제로 바라보게 만든다. HUVM과 memHarvester는 neighbor GPU의 작고 동적인 spare memory를 fast victim cache로 활용하고, eviction/fetch/prefetch path를 병렬화하여 host memory latency를 critical path에서 숨길 수 있음을 보였다. 따라서 이 논문은 GPU UVM, multi-GPU memory management, shared GPU cluster efficiency를 다룰 때 기억할 만한 design point를 제공한다.

[[분류: USENIX ATC]]
[[분류: GPU Unified Virtual Memory]]

GPU Unified Virtual Memory

2026-06-18T04:32:45Z

Noribot: Nori: update draft_gpu_uvm.md

'''GPU Unified Virtual Memory'''는 [[GPU]]와 [[CPU]]가 하나의 virtual address space를 공유하고, data placement와 movement를 runtime/driver가 처리하도록 하는 memory abstraction이다. CUDA 문맥에서는 보통 UVM이라고 부르며, application은 `cudaMallocManaged` 같은 API로 managed memory를 할당하고 CPU/GPU 양쪽에서 같은 pointer를 사용할 수 있다.

핵심은 programmer가 explicit `cudaMemcpy`로 data를 옮기는 대신, GPU가 page를 접근하는 시점에 UVM driver가 필요한 page를 GPU memory로 migrate하거나 host memory에 남겨 둔 채 접근하게 만드는 것이다. 이 abstraction은 GPU memory보다 큰 working set을 실행할 수 있게 해 주지만, page fault handling, migration, prefetching, eviction policy가 성능을 크게 좌우한다.

== 개요 ==
GPU UVM이 제공하려는 것은 두 가지이다.

# '''Pointer sharing''': CPU와 GPU가 같은 virtual address를 통해 같은 data object를 참조할 수 있다.
# '''Demand-driven placement''': data가 처음부터 어느 memory에 있어야 하는지 programmer가 모두 결정하지 않아도, runtime이 접근에 맞춰 page를 이동시킨다.

이 구조 덕분에 GPU application은 GPU memory capacity보다 큰 data set을 host memory까지 활용해 실행할 수 있다. 이를 memory oversubscription이라고 부른다. 그러나 UVM은 "편한 abstraction"과 "빠른 execution"을 자동으로 동시에 보장하지 않는다. GPU가 host memory에 있는 page를 접근하면 fault가 발생하고, 이 fault를 처리하는 동안 GPU execution이 stall될 수 있다.

== 기본 동작 ==
; Managed allocation
: Application은 `cudaMallocManaged` 같은 API로 managed memory object를 할당한다. 이 object는 CPU와 GPU가 모두 접근할 수 있는 virtual address range를 갖는다.

; GPU memory access
: GPU kernel이 managed memory를 load/store할 때 GPU-side MMU가 page table을 확인한다. 해당 page가 GPU memory에 resident하지 않으면 far-fault가 발생한다.

; Far-fault handling
: Far-fault는 GPU memory에 없는 page를 GPU가 접근했을 때 발생하는 fault이다. UVM driver는 fault 정보를 받아 host memory에 있던 page를 GPU memory로 copy하고, page table과 residency metadata를 갱신한 뒤 stalled warp/kernel execution을 재개시킨다.

; Page migration
: UVM은 page를 CPU memory와 GPU memory 사이에서 migrate한다. 논문들에서는 4 KB page fault, 2 MB [[VABlock]] 또는 chunk, 64 KB migration block 같은 여러 granularity가 함께 등장한다. Granularity가 크면 fault 수는 줄일 수 있지만, sparse access에서는 불필요한 data까지 GPU memory를 차지할 수 있다.

; Prefetching
: Fault가 난 page만 옮기면 후속 fault가 계속 발생할 수 있다. 그래서 UVM은 access locality를 가정하고 주변 page를 미리 가져오는 prefetching을 사용한다. NVIDIA UVM의 대표적인 prefetcher로 [[TBNp]]가 있다.

; Eviction
: GPU memory가 부족하면 기존 resident page나 block을 host memory로 evict해야 한다. Eviction policy가 실제 GPU-side reuse를 잘 반영하지 못하면, 방금 evict한 page가 곧 다시 fault를 내며 page thrashing이 발생한다.

== 왜 중요한가 ==
; GPU memory capacity wall
: Deep learning, graph analytics, scientific computing workload는 GPU memory보다 큰 working set을 자주 만든다. UVM은 application을 다시 짜지 않고도 host memory를 capacity extension처럼 사용할 수 있게 해 준다.

; Programmability
: Explicit copy 기반 programming에서는 programmer가 data lifetime, access phase, CPU/GPU ownership을 세밀하게 관리해야 한다. UVM은 같은 pointer를 CPU와 GPU가 공유하게 하므로 code complexity를 낮춘다.

; Performance portability의 유혹
: UVM은 hardware generation, memory size, interconnect bandwidth가 달라도 같은 programming interface를 제공한다. 하지만 실제 성능은 page fault latency, PCIe/NVLink/C2C bandwidth, prefetch/eviction policy에 강하게 의존한다.

== 주요 병목 ==
; Far-fault latency
: GPU가 resident하지 않은 page를 접근하면 fault handling이 critical path에 들어간다. Fault batching이나 pipelining이 없으면 page migration과 page table update 비용이 kernel execution을 크게 늦춘다.

; Coarse-grained migration
: Fault는 작은 page 단위로 발생하더라도 allocation, prefetch, eviction은 더 큰 block/chunk 단위로 일어날 수 있다. Dense access에서는 유리하지만 sparse access에서는 working set을 실제보다 크게 부풀린다.

; PCIe/interconnect bandwidth
: Discrete GPU 환경에서는 CPU memory와 GPU memory 사이의 data movement가 PCIe 또는 NVLink 같은 interconnect를 지난다. GPU local memory bandwidth에 비해 host-device path가 느리면 migration 자체가 병목이 된다.

; Page thrashing
: Oversubscription 상황에서 GPU memory capacity가 부족하면 page가 evict되고 다시 migrate되는 cycle이 생긴다. 잘못된 prefetching은 유용하지 않은 page를 GPU memory에 올려 thrashing을 더 악화시킬 수 있다.

; Driver-observable signal의 한계
: UVM driver는 page fault stream과 driver metadata는 볼 수 있지만, GPU kernel의 실제 data object별 access sequence를 직접 보기 어렵다. 이 한계 때문에 access pattern oblivious policy가 생긴다.

== 대표 정책 ==
; Demand paging
: Page가 실제로 접근될 때 migrate한다. 불필요한 migration은 줄지만, 첫 access마다 fault latency가 발생한다.

; Prefetching
: 앞으로 접근될 가능성이 높은 page를 미리 GPU memory로 가져온다. [[TBNp]]는 tree metadata를 이용해 neighboring page를 prefetch하는 NVIDIA UVM 계열의 대표적인 mechanism이다.

; Access counter-based migration
: GPU가 host memory page를 remote access하도록 두고, access count가 threshold를 넘으면 GPU memory로 migrate하는 방식이다. Sparse access에서 불필요한 migration을 줄일 수 있지만, threshold가 workload나 memory pressure에 맞지 않으면 성능이 흔들린다.

; Zero-copy
: Page를 host memory에 두고 GPU가 remote access한다. Migration을 피할 수 있으므로 sparse access나 severe oversubscription에서는 유리할 수 있지만, GPU local memory보다 bandwidth/latency가 불리하다.

; Placement/eviction policy
: 어떤 page 또는 VABlock을 GPU memory에 둘지, 어떤 data를 host memory로 내보낼지를 결정한다. ARIADNE는 VABlock의 Sharing Degree를 이용해 dense region은 GPU memory에 두고 sparse region은 Zero-copy로 두는 식의 placement policy를 제안한다.

== 정리 ==
GPU Unified Virtual Memory는 GPU programming을 쉽게 만들고 GPU memory capacity를 host memory까지 확장하는 강력한 abstraction이다. 그러나 성능 관점에서는 page fault, migration granularity, prefetching, eviction, Zero-copy placement가 얽힌 복잡한 memory management problem이다. 기억할 점은 UVM이 단순한 "자동 memcpy"가 아니라는 것이다. UVM은 GPU memory와 host memory 사이의 dynamic placement system이고, 좋은 UVM policy는 access pattern, memory pressure, interconnect cost, GPU execution model을 함께 고려해야 한다.

[[분류: GPU Unified Virtual Memory]]

Tree-based Neighboring Prefetcher

2026-06-18T03:06:56Z

Noribot: Nori: update draft_tbnp.md

'''TBNp'''는 '''Tree-based Neighboring Prefetcher'''의 약자로, [[NVIDIA UVM]]에서 [[Unified Virtual Memory]] page fault overhead를 줄이기 위해 사용하는 tree 기반 neighboring prefetcher이다. GPU가 host memory에 있는 UVM page를 접근하면 far-fault가 발생하고, UVM driver가 해당 page를 GPU memory로 migrate한다. TBNp는 이때 fault가 난 page만 가져오는 것이 아니라, 같은 tree 안의 주변 page도 함께 가져와 이후 far-fault를 줄이려는 mechanism이다.

== 개요 ==
[[GPU]] [[Unified Virtual Memory]]에서는 GPU memory보다 큰 working set을 host memory까지 확장해 실행할 수 있다. 하지만 GPU가 host memory에 있는 page를 처음 접근하면 far-fault handling이 필요하고, 이 과정은 GPU execution의 critical path에 들어간다. TBNp는 far-fault가 반복해서 발생하는 것을 줄이기 위해 spatial locality를 가정하고 neighboring page를 미리 migrate한다.

TBNp는 다음 구조를 갖는다.
# UVM virtual address space를 2 MB [[VABlock]] 단위로 나눈다.
# 각 2 MB VABlock을 5-level full binary tree로 관리한다.
# leaf node는 64 KB basic block이며, TBNp의 최소 migration unit처럼 동작한다.
# non-leaf node는 subtree 전체 크기인 <math>N_{total}</math>과 이미 GPU memory로 migrate된 child node 크기인 <math>N_{migrated}</math>를 metadata로 가진다.
# subtree에서 migrate된 비율이 50%를 넘으면, locality가 높다고 보고 남은 child node들을 proactive하게 prefetch한다.

즉 TBNp는 page fault stream을 보고 "이 근처도 곧 쓸 것이다"라고 추정한다. 이 추정을 tree metadata로 관리하기 때문에, 단순히 fixed-size window를 prefetch하는 방식보다 hierarchy별 locality를 표현할 수 있다.

== 동작 방식 ==
; Far-fault 발생
: GPU가 resident하지 않은 UVM page를 접근하면 far-fault가 발생한다. UVM driver는 fault가 난 page가 포함된 TBNp leaf node를 찾는다.

; Leaf-node migration
: conventional TBNp에서는 leaf node가 64 KB이므로, 4 KB page 하나에서 fault가 나도 같은 64 KB leaf node 안의 나머지 page까지 GPU memory로 migrate한다. 이 단계만으로도 같은 leaf 안의 후속 page fault를 줄일 수 있다.

; Tree metadata update
: migration이 일어나면 leaf에서 root 방향으로 <math>N_{migrated}</math>가 갱신된다. 각 non-leaf node는 자기 subtree에서 얼마나 많은 child node가 GPU memory에 있는지를 알 수 있다.

; Proactive prefetch
: 어떤 non-leaf node에서 <math>N_{migrated}</math>가 <math>N_{total}</math>의 50%를 넘으면, TBNp는 그 subtree의 나머지 leaf node도 GPU memory로 prefetch한다. 예를 들어 같은 2 MB VABlock의 앞쪽 leaf들이 이미 많이 migrate되었다면, 나머지 인접 leaf도 곧 접근될 것이라고 본다.

; Far-fault amortization
: TBNp의 목적은 한 번의 far-fault handling 동안 여러 page를 함께 가져와, 이후 발생할 수 있는 여러 on-demand far-fault를 제거하는 것이다. GPU workload가 sequential 또는 dense spatial access를 보이면 이 전략이 잘 맞는다.

=== 왜 tree를 쓰는가 ===
GPU memory access는 CPU prefetcher가 가정하는 instruction-local cache miss stream과 다르게, 많은 warp와 SM에서 bursty하게 발생할 수 있다. 개별 load instruction 단위로 fine-grained access pattern을 추적하면 metadata와 driver communication 비용이 커진다. 반대로 TBNp는 page migration이라는 coarse-grained UVM event를 tree 구조로 묶어, batch 단위 locality와 prefetch aggressiveness를 비교적 간단하게 조절한다.

Tree-based prefetching이 GPU의 bursty memory behavior에 잘 맞고, prefetch aggressiveness를 hierarchical하게 조절할 수 있기 때문에 TBNp가 최근 NVIDIA GPU UVM에서 중요하게 사용되고 있다.

== 한계 ==
; Homogeneous configuration
: conventional TBNp는 모든 application, kernel, data object에 같은 tree configuration을 적용한다. Forest 논문에서 다룬 baseline은 2 MB tree와 64 KB leaf node를 사용한다. 그러나 실제 workload에서는 어떤 object는 linear streaming이고, 어떤 object는 sparse irregular access를 보인다. 하나의 tree size와 leaf size가 모든 object에 최적일 수 없다.

; Unnecessary migration
: TBNp는 locality를 가정하고 neighboring page를 미리 가져온다. 하지만 sparse access에서는 64 KB leaf 안에서도 실제로 쓰는 4 KB page가 일부뿐일 수 있다. 이 경우 migration된 page가 eviction 전까지 한 번도 접근되지 않을 수 있고, GPU memory capacity를 불필요하게 점유한다.

; Oversubscription에서 memory pressure 악화
: GPU memory가 oversubscribed된 상황에서는 prefetch가 항상 좋은 것이 아니다. 잘못 prefetch된 page가 GPU memory를 차지하면 유용한 page가 밀려나고, 이후 다시 migrate되면서 page thrashing이 생긴다. Forest 논문은 일부 benchmark에서 migrated page의 5%에서 48%가 eviction 전까지 전혀 접근되지 않았다고 보고한다.

; Eviction policy와의 상호작용
: conventional UVM eviction은 GPU-side actual access recency가 아니라 far-fault event history에 의존한다. 따라서 TBNp가 불필요한 migration을 만들면 eviction candidate도 왜곡될 수 있다. Prefetcher와 eviction policy는 독립된 문제가 아니라 같은 GPU memory capacity를 두고 상호작용한다.

; Access pattern obliviousness
: TBNp는 page fault와 migration state는 보지만, data object의 semantic이나 GPU-side access sequence를 직접 이해하지 않는다. 그래서 같은 application 안에서 kernel마다 object access pattern이 달라져도 동일한 policy를 적용한다.

[[분류: GPU Unified Virtual Memory]]

ARIADNE: Adaptive UVM Management for Efficient GPU Memory Oversubscription

2026-06-17T09:15:13Z

Noribot: Nori: update draft.md

{{Paper
|title=ARIADNE: Adaptive UVM Management for Efficient GPU Memory Oversubscription
|author=Hyunkyun Shin, Seongtae Bang, Hyungwon Park, Daehoon Kim
|conference=IEEE International Symposium on High Performance Computer Architecture (HPCA)
|year=2026
|doi=10.1109/HPCA68181.2026.11408564
}}

== 개요 ==
이 논문은 [[GPU]] [[Unified Virtual Memory]]에서 memory oversubscription이 발생할 때 page fault 처리 지연과 thrashing이 왜 급격한 성능 저하로 이어지는지 분석하고, runtime-only UVM driver management인 ARIADNE로 memory region placement를 동적으로 바꾸어 이를 완화할 수 있음을 보였다. 기존 [[NVIDIA UVM]]의 fault handling path와 page placement policy를 바꾸되, hardware, compiler, application code를 수정하지 않고도 최적화할 수 있는 transparency를 제공한다.

== Motivation ==
; Page fault handling overhead가 큼
: Host memory에 있는 UVM page를 GPU memory로 migration하기 위한 fault handling 비용이 크다. NVIDIA의 access counter-based migration(AC)은 page를 initially Zero-copy state에 두고, access count가 static threshold를 넘는 page만 GPU로 migrate하여 불필요한 migration을 줄이려 한다. 그러나 threshold가 workload의 access pattern이나 현재 memory pressure에 적응하지 못하기 때문에, no-oversubscription 상황에서는 Zero-copy overhead가 커지고 high-oversubscription 상황에서는 thrashing을 완전히 막지 못한다.

; oversubscription 상황에서 UVM이 매우 쉽게 thrashing에 빠짐
: 기본 UVM은 4 KB page fault를 받더라도 GPU physical memory allocation은 보통 2 MB VABlock/chunk 단위로 수행한다. Sparse access에서는 실제로 필요한 page보다 큰 chunk가 GPU memory를 점유하면서 WCSS가 부풀고, last-fault-time 기반 victim selection은 아직 active한 VABlock을 evict한 뒤 곧바로 다시 refetch하게 만들 수 있다. 이 coarse-grained allocation과 단순 eviction policy의 조합이 repeated eviction-refetch cycle을 만든다.

기존 대응은 크게 prefetching, access counter 기반 migration, Zero-copy placement로 나뉜다. 그러나 prefetching은 thrashing 자체를 제거하지 못하고, access counter 방식은 static threshold에 의존하며, compiler/hardware 기반 Zero-copy 선택은 UVM의 장점인 binary transparency와 hardware portability를 약화시킨다.

== Main Idea ==
; Design Principle #1. Reduce UVM fault-handling latency through pipelined execution
: ARIADNE는 UVM fault handling의 latency bottleneck이 Copy/Eviction보다 Populate에 있다는 관찰을 사용한다. Populate, Copy, Eviction을 VABlock 간 pipeline으로 분리하여 long-latency Populate를 다른 VABlock의 Copy/Eviction과 overlap시킨다. 즉 ARIADNE는 placement policy와 fault handling pipeline을 함께 바꾸어 thrashing과 migration latency를 동시에 줄이려 한다.

; Design Principle #2. Sharing Degree: a runtime access-pattern metric leveraging thread-level information
: GPU thread execution architecture가 VABlock-level spatial locality의 runtime signal을 남긴다는 관찰을 이용하는 것이다. GPU kernel에서 thread block은 SM에 배치되고, 인접 SM들은 uTLB를 공유한다. 어떤 VABlock을 여러 uTLB/SM이 동시에 faulting한다면, 서로 다른 thread block이 그 VABlock 내부의 여러 page를 접근하고 있을 가능성이 높다. 반대로 unique uTLB가 적으면 일부 thread group만 sparse하게 접근하는 영역일 가능성이 높다. 이를 반영하는 Sharing Degree라는 runtime metric을 정의하였다. UVM driver는 kernel source나 per-thread address expression을 보지 못하지만, page fault record의 source uTLB ID는 볼 수 있다. ARIADNE는 최근 fault를 발생시킨 unique uTLB 수를 VABlock별로 추적하여, 해당 VABlock을 동시에 접근하는 SM/thread group의 수를 근사한다. 이 값이 높으면 VABlock 내부의 page utilization이 높고 dense access일 가능성이 크다는 것이 논문의 핵심 관찰이다.

; Design Principle #3. Managing memory region placement between GPU memory and Zero-copy based on runtime memory access characteristics
: 특히 migration은 2 MB chunk 단위이고 Zero-copy는 128 B cache line 단위이므로, dense VABlock은 GPU memory에 두는 것이 유리하고 sparse VABlock은 host memory의 Zero-copy로 두는 것이 유리하다는 placement framing을 제시한다. ARIADNE는 Sharing Degree를 이용해 VABlock을 GPU memory와 Zero-copy 사이에서 동적으로 배치한다. Dense VABlock은 2 MB chunk를 GPU에 올려도 내부 활용도가 높으므로 migration이 유리하다. Sparse VABlock은 chunk를 통째로 점유하면 WCSS를 부풀리므로, evicted 이후 재접근되더라도 곧장 GPU로 다시 올리지 않고 일시적으로 Zero-copy로 둔다.

== Background ==
; VABlock과 chunk
: NVIDIA UVM은 virtual address range를 VABlock으로 관리하고, GPU physical memory는 chunk 단위로 할당한다. 논문 기준으로 보통 VABlock과 chunk는 각각 2 MB이다. Page fault는 4 KB 단위로 발생하지만 allocation과 eviction은 더 coarse-grained하게 일어난다.

; Populate, Eviction, Copy
: UVM fault handling은 GPU chunk를 준비하는 Populate, victim chunk를 host로 내보내는 Eviction, host page를 GPU memory로 복사하는 Copy로 구성된다. 논문은 Populate latency가 Copy/Eviction의 거의 두 배이며, fault batch 전체 latency를 지배한다고 분석한다.

; Zero-copy
: Zero-copy는 page를 host memory에 pin한 뒤 GPU가 remote access하도록 하는 방식이다. GPU memory bandwidth보다 느리지만, sparse VABlock을 2 MB chunk로 migrate하지 않아도 되므로 oversubscription thrashing을 줄일 수 있다.

== Design ==
; Working Chunk Set Size (WCSS) estimation
: ARIADNE는 workload가 실제로 요구하는 GPU chunk 수를 WCSS로 추적한다. 단순히 GPU-resident VABlock만 세면 방금 evict되었지만 곧 재접근될 active VABlock을 놓칠 수 있으므로, GPU resident VABlock, Zero-copy VABlock, re-access 가능성이 높은 recently evicted VABlock을 함께 포함한다. 초기 정보가 부족할 때는 GPU-wide average Sharing Degree가 낮으면 sparse workload로 보고 최근 evicted VABlock을 500 ms 동안 WCSS에 보존한다.

; Sharing Degree measurement
: 각 VABlock에 대해 최근 16개 page fault의 source uTLB ID를 circular queue에 저장하고, unique uTLB ID 개수를 Sharing Degree로 계산한다. Source uTLB ID는 UVM driver의 fault record에서 얻을 수 있으므로 hardware나 compiler modification이 필요 없다. 논문은 ATAX/GEMV 같은 sparse workload는 Sharing Degree가 주로 1이고, GEMM/HEL 같은 dense workload는 4보다 큰 값이 많으며, BFS/NW/XSB처럼 mixed pattern인 workload는 VABlock별로 다른 값을 보인다고 제시하였다. 이 Sharing Degree는 VABlock을 접근하는 SM/thread group의 다양성을 통해 spatial locality를 근사한다. 따라서 ARIADNE에서는 '''Sharing Degree로 VABlock이 sparse한지 dense한지 판단'''하는 것이 placement decision의 핵심이다.

; Pipelined VABlock fault handling
: 기본 UVM은 Populate, Eviction, Copy를 monolithic sequential routine으로 처리한다. ARIADNE는 Populate를 Copy와 분리하고, 다른 VABlock의 Populate/Copy/Eviction을 병렬로 실행한다. Populate는 GPU chunk allocation과 metadata 준비만 끝낸 뒤 Copy worker에 넘기고, Eviction은 free chunk가 부족해진 뒤 반응적으로 실행되는 대신 dedicated thread에서 proactive하게 실행된다. 이 구조는 Populate의 긴 latency를 다른 작업 뒤에 숨기는 것이 목적이다.

; Sharing Degree-aware eviction
: ARIADNE는 fault recency만 보는 eviction queue 대신 Sharing Degree를 반영한 priority key를 사용한다. 논문이 제시한 형태는 <math>last fault time + (SD Weight * Sharing Degree / Nfault_history)</math>이며, 구현에서는 fault history 길이 16과 SD Weight 100 us를 사용한다. Sharing Degree가 높은 VABlock은 GPU memory에 남길 가치가 높다고 보고 eviction priority를 낮춘다.

; Dynamic Zero-copy
: Memory demand가 GPU capacity를 넘으면 ARIADNE는 evicted 뒤 재접근된 VABlock을 즉시 GPU로 refetch하지 않고 100 ms 동안 Zero-copy state로 유지한다. 이 second-chance policy는 방금 쫓겨난 sparse VABlock이 다시 fault를 내며 eviction-refetch cycle을 만드는 것을 줄인다. 100 ms가 지나도록 재사용되지 않으면 Zero-copy state를 해제하고, 이후에도 계속 재사용되면 GPU memory로 promote될 수 있다.

; Dynamic prefetching
: Copy 단계에서는 기본 UVM prefetcher가 선택한 page를 복사하되, GPU memory 여유가 충분하거나 VABlock의 Sharing Degree가 threshold보다 높으면 VABlock 전체를 적극적으로 copy한다. 논문 구현의 Sharing Degree threshold는 3이다. 이 정책은 dense VABlock에 대해서는 future fault를 줄이고, sparse VABlock에 대해서는 불필요한 2 MB migration을 피하려는 tradeoff를 갖는다.

== Result ==
평가는 NVIDIA RTX A5000, AMD Ryzen 7700X, PCIe 4.0 x16, 64 GB DDR5, Linux 6.0, NVIDIA open-source kernel driver 535.86 환경에서 수행되었다. Workload는 Rodinia, Polybench, HeCBench, XSBench에서 가져온 10개 benchmark(2DC, ATAX, BICG, GEMM, GEMV, MVT, XSB, BFS, HEL, NW)이며, 각 benchmark의 memory footprint는 4 GB로 맞추었다. Oversubscription ratio는 no oversubscription, 130%, 175%, 300%를 사용한다.

Baseline UVM은 oversubscription에서 급격히 무너진다. 논문은 baseline UVM의 geomean execution time이 200% oversubscription에서 33.2배, 300%에서 60.7배까지 증가한다고 보고한다. AC(access counter-based migration)는 Zero-copy를 활용해 thrashing을 줄이지만 no oversubscription에서는 평균 1.3배, 최대 2.5배 느려지고, 210%와 300% oversubscription에서는 no oversubscription 대비 각각 4.7배와 8.9배 느려진다.

ARIADNE는 AC 대비 130%, 175%, 300% oversubscription에서 각각 1.9배, 2.3배, 4.0배 geomean speedup을 보인다. SUV 대비로는 같은 oversubscription ratio에서 각각 1.9배, 5.0배, 4.8배 speedup을 보인다. 또한 no-oversubscription 대비 runtime 증가는 130%, 175%, 300%에서 각각 1.6배, 1.8배, 2.3배로 보고되어, 논문은 near-linear degradation이라고 해석한다.

Mechanism-level evidence도 제시된다. Pipelined VABlock fault handling은 10개 benchmark에서 VABlock fault handling latency를 평균 17%, 최대 48%(BFS) 줄인다. Dynamic VABlock placement는 175% oversubscription에서 AC 대비 PCIe traffic을 평균 51% 수준으로 낮춘다. Breakdown에서는 Sharing Degree 기반 placement/eviction이 핵심 성능 요인이고, pipelining이 oversubscription 상황에서 정책 overhead를 숨기는 역할을 한다고 해석된다.

LLM inference 실험도 포함된다. Llama3.1 70B inference(input token length 2048)에서 ARIADNE는 AC 대비 Decode phase 4.2배, Prefill phase 1.6배 speedup을 보인다. Decode가 GEMV 중심이기 때문에 GEMV benchmark에서 큰 이득을 보인 결과와 일관된다고 설명한다.

Overhead는 작다고 보고된다. Additional metadata는 VABlock당 70 B 미만, GPU당 100 B 미만이며, 16 GB application 기준 약 560 KB이다. Sharing Degree/WCSS tracking과 Zero-copy process의 추가 latency는 최대 100 ns로, 단일 VABlock fault handling 약 20 us에 비해 무시 가능하다고 주장한다.

== Contribution ==
# UVM oversubscription에서 성능 저하의 원인을 Populate-dominated fault handling latency, 2 MB chunk granularity에 따른 WCSS amplification, last-fault-time eviction에 따른 thrashing으로 정리하였다.
# GPU thread execution architecture와 source uTLB ID를 이용해 VABlock-level spatial locality를 runtime에서 추정하는 Sharing Degree metric을 제안하였다.
# Sharing Degree와 WCSS를 기반으로 VABlock을 GPU memory와 Zero-copy 사이에서 동적으로 배치하는 runtime-only UVM management framework ARIADNE를 설계하였다.
# Populate, Copy, Eviction을 VABlock 간 pipeline으로 분리하여 UVM fault handling latency를 줄이는 driver-level execution structure를 구현하였다.
# NVIDIA open-source UVM driver 내부 수정 약 1600 LOC만으로 구현하고, hardware/compiler/application modification 없이 executable 또는 closed-source UVM application에 적용 가능한 설계를 보였다.
# 10개 GPU benchmark와 Llama3.1 70B inference에서 AC, SUV, baseline UVM 대비 성능 향상과 oversubscription scaling을 실험적으로 제시하였다.

== Criticisms ==
# Sharing Degree는 source uTLB ID를 thread group locality의 proxy로 사용하는 metric이다. 논문은 benchmark에서 utilization과의 상관을 보이지만, uTLB sharing topology, SM scheduling, phase behavior가 다른 GPU generation에서 threshold 3, history length 16, SD Weight 100 us, Zero-copy pin time 100 ms가 항상 좋은지는 제한적이다. 따라서 ARIADNE의 parameter choice가 특정 실험 환경에 과도하게 맞춰진 것은 아닌지 추가 검증할 필요가 있다.
# 전체적으로 ARIADNE는 여러 driver-level optimization을 조합한 system으로 보인다. 따라서 각 mechanism이 서로 다른 workload와 memory pressure에서 독립적으로 얼마나 기여하는지, 그리고 parameter가 바뀌어도 같은 design principle이 유지되는지를 더 강하게 분리해 보여주면 설득력이 커질 수 있다. 다만 Sharing Degree 자체는 UVM driver가 관찰 가능한 정보만으로 VABlock locality를 추정한다는 점에서 명확한 design point를 제공한다.
# no-oversubscription 또는 memory pressure가 낮은 일부 case에서는 SUV가 ARIADNE보다 빠르다. 논문은 SUV의 compile-time range prefetching이 ARIADNE의 2 MB VABlock granularity prefetch보다 큰 data range를 미리 가져올 수 있기 때문이라고 설명한다. 즉 ARIADNE는 transparency를 얻는 대신 static program knowledge를 활용한 aggressive prefetch opportunity를 일부 잃는다.

== [[Conclusion]] ==
이 연구는 GPU UVM oversubscription을 단순 page fault overhead가 아니라 VABlock placement와 GPU thread-level locality를 함께 다루어야 하는 문제로 바라보게 만든다. ARIADNE는 source uTLB 기반 Sharing Degree, WCSS estimation, Sharing Degree-aware eviction, transient Zero-copy, pipelined fault handling을 결합하여 hardware/compiler/application 수정 없이 oversubscription 성능을 개선할 수 있음을 보였다.

[[분류: 시스템 논문]]
[[분류: IEEE HPCA]]
[[분류: GPU Unified Virtual Memory]]

ARIADNE: Adaptive UVM Management for Efficient GPU Memory Oversubscription

2026-06-17T09:02:03Z

Noribot: Nori: update draft.md

{{Paper
|title=ARIADNE: Adaptive UVM Management for Efficient GPU Memory Oversubscription
|author=Hyunkyun Shin, Seongtae Bang, Hyungwon Park, Daehoon Kim
|conference=IEEE International Symposium on High Performance Computer Architecture (HPCA)
|year=2026
|doi=10.1109/HPCA68181.2026.11408564
}}

== 개요 ==
이 논문은 [[GPU]] [[Unified Virtual Memory]]에서 memory oversubscription이 발생할 때 page fault 처리 지연과 thrashing이 왜 급격한 성능 저하로 이어지는지 분석하고, runtime-only UVM driver management인 ARIADNE로 memory region placement를 동적으로 바꾸어 이를 완화할 수 있음을 보였다. 기존 [[NVIDIA UVM]]의 fault handling path와 page placement policy를 바꾸되, hardware, compiler, application code를 수정하지 않고도 최적화 할수 있는 Transparency를 제공한다.

== Motivation ==
; Pae fault handling overhead가 큼
: CPU메모리를 GPU로 Migration하기 위한 비용이 크다. 이를 줄이기 위해서 기존에는 access counter값이 특정 수치 이상이면, zero-copy를 수행하는 (Nvdia AC)기법을 활용했지만, 특정 workload에서는 오히려 성능이 감소하는 문제가 발생하였다.

; oversubscription 상황에서 UVM이 매우 쉽게 thrashing에 빠짐
: 기본 UVM은 4 KB page fault를 받더라도 GPU physical memory allocation은 보통 2 MB VABlock/chunk 단위로 수행한다. 따라서 CPU page evction policy와 GPU page evction policy가 일치하지 않는 문제가 발생한다.

기존 대응은 크게 prefetching, access counter 기반 migration, Zero-copy placement로 나뉜다. 그러나 prefetching은 thrashing 자체를 제거하지 못하고, access counter 방식은 static threshold에 의존하며, compiler/hardware 기반 Zero-copy 선택은 UVM의 장점인 binary transparency와 hardware portability를 약화시킨다.

== Main Idea ==
; Design Principle #1. Reduce UVM fault-handling latency through pipelined execution
: ARIADNE는 UVM fault handling의 latency bottleneck이 Copy/Eviction보다 Populate에 있다는 관찰을 사용한다. Populate, Copy, Eviction을 VABlock 간 pipeline으로 분리하여 long-latency Populate를 다른 VABlock의 Copy/Eviction과 overlap시킨다. 즉 ARIADNE는 placement policy와 fault handling pipeline을 함께 바꾸어 thrashing과 migration latency를 동시에 줄이려 한다.

; Design Principle #2. Sharing Degree: a runtime access-pattern metric leveraging thread-level information
: GPU thread execution architecture가 VABlock-level spatial locality의 runtime signal을 남긴다는 관찰을 이용하는 것이다. GPU kernel에서 thread block은 SM에 배치되고, 인접 SM들은 uTLB를 공유한다. 어떤 VABlock을 여러 uTLB/SM이 동시에 faulting한다면, 서로 다른 thread block이 그 VABlock 내부의 여러 page를 접근하고 있을 가능성이 높다. 반대로 unique uTLB가 적으면 일부 thread group만 sparse하게 접근하는 영역일 가능성이 높다. 이를 반영하는 Sharing Degree라는 runtime metric을 정의하였다. UVM driver는 kernel source나 per-thread address expression을 보지 못하지만, page fault record의 source uTLB ID는 볼 수 있다. ARIADNE는 최근 fault를 발생시킨 unique uTLB 수를 VABlock별로 추적하여, 해당 VABlock을 동시에 접근하는 SM/thread group의 수를 근사한다. 이 값이 높으면 VABlock 내부의 page utilization이 높고 dense access일 가능성이 크다는 것이 논문의 핵심 관찰이다.

; Design Principle #3. Managing CPU memory region placement between GPU memory and Zero-copy based on runtime memory access characteristics
특히 migration은 2 MB chunk 단위이고 Zero-copy는 128 B cache line 단위이므로, dense VABlock은 GPU memory에 두는 것이 유리하고 sparse VABlock은 host memory의 Zero-copy로 두는 것이 유리하다는 placement framing을 제시한다. ARIADNE는 Sharing Degree를 이용해 VABlock을 GPU memory와 Zero-copy 사이에서 동적으로 배치한다. Dense VABlock은 2 MB chunk를 GPU에 올려도 내부 활용도가 높으므로 migration이 유리하다. Sparse VABlock은 chunk를 통째로 점유하면 WCSS를 부풀리므로, evicted 이후 재접근되더라도 곧장 GPU로 다시 올리지 않고 일시적으로 Zero-copy로 둔다.

== Background ==
; VABlock과 chunk
: NVIDIA UVM은 virtual address range를 VABlock으로 관리하고, GPU physical memory는 chunk 단위로 할당한다. 논문 기준으로 보통 VABlock과 chunk는 각각 2 MB이다. Page fault는 4 KB 단위로 발생하지만 allocation과 eviction은 더 coarse-grained하게 일어난다.

; Populate, Eviction, Copy
: UVM fault handling은 GPU chunk를 준비하는 Populate, victim chunk를 host로 내보내는 Eviction, host page를 GPU memory로 복사하는 Copy로 구성된다. 논문은 Populate latency가 Copy/Eviction의 거의 두 배이며, fault batch 전체 latency를 지배한다고 분석한다.

; Zero-copy
: Zero-copy는 page를 host memory에 pin한 뒤 GPU가 remote access하도록 하는 방식이다. GPU memory bandwidth보다 느리지만, sparse VABlock을 2 MB chunk로 migrate하지 않아도 되므로 oversubscription thrashing을 줄일 수 있다.

== Design ==
; Working Chunk Set Size (WCSS) estimation
: ARIADNE는 workload가 실제로 요구하는 GPU chunk 수를 WCSS로 추적한다. 단순히 GPU-resident VABlock만 세면 방금 evict되었지만 곧 재접근될 active VABlock을 놓칠 수 있으므로, GPU resident VABlock, Zero-copy VABlock, re-access 가능성이 높은 recently evicted VABlock을 함께 포함한다. 초기 정보가 부족할 때는 GPU-wide average Sharing Degree가 낮으면 sparse workload로 보고 최근 evicted VABlock을 500 ms 동안 WCSS에 보존한다.

; Sharing Degree measurement
: 각 VABlock에 대해 최근 16개 page fault의 source uTLB ID를 circular queue에 저장하고, unique uTLB ID 개수를 Sharing Degree로 계산한다. Source uTLB ID는 UVM driver의 fault record에서 얻을 수 있으므로 hardware나 compiler modification이 필요 없다. 논문은 ATAX/GEMV 같은 sparse workload는 Sharing Degree가 주로 1이고, GEMM/HEL 같은 dense workload는 4보다 큰 값이 많으며, BFS/NW/XSB처럼 mixed pattern인 workload는 VABlock별로 다른 값을 보인다고 제시하였다. 이 Sharing Degree는 SM들에 접근하는 메모리의 spatial locality에 대한 정보를 효과적으로 보여준다고 주장한다. 따라서 '''Sparse memory의 값으로 sparse인지 아니면 dense한지 판단'''하는 것이 중요하다고 주장하였다.

; Pipelined VABlock fault handling
: 기본 UVM은 Populate, Eviction, Copy를 monolithic sequential routine으로 처리한다. ARIADNE는 Populate를 Copy와 분리하고, 다른 VABlock의 Populate/Copy/Eviction을 병렬로 실행한다. Populate는 GPU chunk allocation과 metadata 준비만 끝낸 뒤 Copy worker에 넘기고, Eviction은 free chunk가 부족해진 뒤 반응적으로 실행되는 대신 dedicated thread에서 proactive하게 실행된다. 이 구조는 Populate의 긴 latency를 다른 작업 뒤에 숨기는 것이 목적이다.

; Sharing Degree-aware eviction
: ARIADNE는 fault recency만 보는 eviction queue 대신 Sharing Degree를 반영한 priority key를 사용한다. 논문이 제시한 형태는 <math>last fault time + (SD Weight * Sharing Degree / Nfault_history)</math>이며, 구현에서는 fault history 길이 16과 SD Weight 100 us를 사용한다. Sharing Degree가 높은 VABlock은 GPU memory에 남길 가치가 높다고 보고 eviction priority를 낮춘다.

; Dynamic Zero-copy
: Memory demand가 GPU capacity를 넘으면 ARIADNE는 evicted 뒤 재접근된 VABlock을 즉시 GPU로 refetch하지 않고 100 ms 동안 Zero-copy state로 유지한다. 이 second-chance policy는 방금 쫓겨난 sparse VABlock이 다시 fault를 내며 eviction-refetch cycle을 만드는 것을 줄인다. 100 ms가 지나도록 재사용되지 않으면 Zero-copy state를 해제하고, 이후에도 계속 재사용되면 GPU memory로 promote될 수 있다.

; Dynamic prefetching
: Copy 단계에서는 기본 UVM prefetcher가 선택한 page를 복사하되, GPU memory 여유가 충분하거나 VABlock의 Sharing Degree가 threshold보다 높으면 VABlock 전체를 적극적으로 copy한다. 논문 구현의 Sharing Degree threshold는 3이다. 이 정책은 dense VABlock에 대해서는 future fault를 줄이고, sparse VABlock에 대해서는 불필요한 2 MB migration을 피하려는 tradeoff를 갖는다.

== Result ==
평가는 NVIDIA RTX A5000, AMD Ryzen 7700X, PCIe 4.0 x16, 64 GB DDR5, Linux 6.0, NVIDIA open-source kernel driver 535.86 환경에서 수행되었다. Workload는 Rodinia, Polybench, HeCBench, XSBench에서 가져온 10개 benchmark(2DC, ATAX, BICG, GEMM, GEMV, MVT, XSB, BFS, HEL, NW)이며, 각 benchmark의 memory footprint는 4 GB로 맞추었다. Oversubscription ratio는 no oversubscription, 130%, 175%, 300%를 사용한다.

Baseline UVM은 oversubscription에서 급격히 무너진다. 논문은 baseline UVM의 geomean execution time이 200% oversubscription에서 33.2배, 300%에서 60.7배까지 증가한다고 보고한다. AC(access counter-based migration)는 Zero-copy를 활용해 thrashing을 줄이지만 no oversubscription에서는 평균 1.3배, 최대 2.5배 느려지고, 210%와 300% oversubscription에서는 no oversubscription 대비 각각 4.7배와 8.9배 느려진다.

ARIADNE는 AC 대비 130%, 175%, 300% oversubscription에서 각각 1.9배, 2.3배, 4.0배 geomean speedup을 보인다. SUV 대비로는 같은 oversubscription ratio에서 각각 1.9배, 5.0배, 4.8배 speedup을 보인다. 또한 no-oversubscription 대비 runtime 증가는 130%, 175%, 300%에서 각각 1.6배, 1.8배, 2.3배로 보고되어, 논문은 near-linear degradation이라고 해석한다.

Mechanism-level evidence도 제시된다. Pipelined VABlock fault handling은 10개 benchmark에서 VABlock fault handling latency를 평균 17%, 최대 48%(BFS) 줄인다. Dynamic VABlock placement는 175% oversubscription에서 AC 대비 PCIe traffic을 평균 51% 수준으로 낮춘다. Breakdown에서는 Sharing Degree 기반 placement/eviction이 핵심 성능 요인이고, pipelining이 oversubscription 상황에서 정책 overhead를 숨기는 역할을 한다고 해석된다.

LLM inference 실험도 포함된다. Llama3.1 70B inference(input token length 2048)에서 ARIADNE는 AC 대비 Decode phase 4.2배, Prefill phase 1.6배 speedup을 보인다. Decode가 GEMV 중심이기 때문에 GEMV benchmark에서 큰 이득을 보인 결과와 일관된다고 설명한다.

Overhead는 작다고 보고된다. Additional metadata는 VABlock당 70 B 미만, GPU당 100 B 미만이며, 16 GB application 기준 약 560 KB이다. Sharing Degree/WCSS tracking과 Zero-copy process의 추가 latency는 최대 100 ns로, 단일 VABlock fault handling 약 20 us에 비해 무시 가능하다고 주장한다.

== Contribution ==
# UVM oversubscription에서 성능 저하의 원인을 Populate-dominated fault handling latency, 2 MB chunk granularity에 따른 WCSS amplification, last-fault-time eviction에 따른 thrashing으로 정리하였다.
# GPU thread execution architecture와 source uTLB ID를 이용해 VABlock-level spatial locality를 runtime에서 추정하는 Sharing Degree metric을 제안하였다.
# Sharing Degree와 WCSS를 기반으로 VABlock을 GPU memory와 Zero-copy 사이에서 동적으로 배치하는 runtime-only UVM management framework ARIADNE를 설계하였다.
# Populate, Copy, Eviction을 VABlock 간 pipeline으로 분리하여 UVM fault handling latency를 줄이는 driver-level execution structure를 구현하였다.
# NVIDIA open-source UVM driver 내부 수정 약 1600 LOC만으로 구현하고, hardware/compiler/application modification 없이 executable 또는 closed-source UVM application에 적용 가능한 설계를 보였다.
# 10개 GPU benchmark와 Llama3.1 70B inference에서 AC, SUV, baseline UVM 대비 성능 향상과 oversubscription scaling을 실험적으로 제시하였다.

== Criticisms ==
# Sharing Degree는 source uTLB ID를 thread group locality의 proxy로 사용하는 metric이다. 논문은 benchmark에서 utilization과의 상관을 보이지만, uTLB sharing topology, SM scheduling, phase behavior가 다른 GPU generation에서 threshold 3, history length 16, SD Weight 100 us, Zero-copy pin time 100 ms가 항상 좋은지는 제한적이다. 약간 fine-tuning되어 있는 값처럼 보이는 경향이 있다.
# 전체적으로 논문이 Collection of fine-tuned optimization으로 보이는 경향이 있다. 그러나 논문에서 제시되는 Sharing Degree와 같은 개념은 좋은 Design point이기 떄문에, Collection-of-optimizations으로 보는 것이 오히려 불필요하게 Critical한 Review일 수 있을 것 같다는 생각도 든다. 이런 Optimization work의 숙명일 수도....
# no-oversubscription 또는 memory pressure가 낮은 일부 case에서는 SUV가 ARIADNE보다 빠르다. 논문은 SUV의 compile-time range prefetching이 ARIADNE의 2 MB VABlock granularity prefetch보다 큰 data range를 미리 가져올 수 있기 때문이라고 설명한다. 즉 ARIADNE는 transparency를 얻는 대신 static program knowledge를 활용한 aggressive prefetch opportunity를 일부 잃는다.

== [[Conclusion]] ==
이 연구는 GPU UVM oversubscription을 단순 page fault overhead가 아니라 VABlock placement와 GPU thread-level locality를 함께 다루어야 하는 문제로 바라보게 만든다. ARIADNE는 source uTLB 기반 Sharing Degree, WCSS estimation, Sharing Degree-aware eviction, transient Zero-copy, pipelined fault handling을 결합하여 hardware/compiler/application 수정 없이 oversubscription 성능을 개선할 수 있음을 보였다.

[[분류: 시스템 논문]]
[[분류: IEEE HPCA]]
[[분류: GPU Unified Virtual Memory]]

ARIADNE: Adaptive UVM Management for Efficient GPU Memory Oversubscription

2026-06-17T09:00:41Z

Noribot: Nori: update draft.md

{{Paper
|title=ARIADNE: Adaptive UVM Management for Efficient GPU Memory Oversubscription
|author=Hyunkyun Shin, Seongtae Bang, Hyungwon Park, Daehoon Kim
|conference=IEEE International Symposium on High Performance Computer Architecture (HPCA)
|year=2026
|doi=10.1109/HPCA68181.2026.11408564
}}

== 개요 ==
이 논문은 [[GPU]] [[Unified Virtual Memory]]에서 memory oversubscription이 발생할 때 page fault 처리 지연과 thrashing이 왜 급격한 성능 저하로 이어지는지 분석하고, runtime-only UVM driver management인 ARIADNE로 memory region placement를 동적으로 바꾸어 이를 완화할 수 있음을 보였다. 기존 [[NVIDIA UVM]]의 fault handling path와 page placement policy를 바꾸되, hardware, compiler, application code를 수정하지 않고도 최적화 할수 있는 Transparency를 제공한다.

== Motivation ==
; Pae fault handling overhead가 큼
: CPU메모리를 GPU로 Migration하기 위한 비용이 크다. 이를 줄이기 위해서 기존에는 access counter값이 특정 수치 이상이면, zero-copy를 수행하는 (Nvdia AC)기법을 활용했지만, 특정 workload에서는 오히려 성능이 감소하는 문제가 발생하였다.

; oversubscription 상황에서 UVM이 매우 쉽게 thrashing에 빠짐
: 기본 UVM은 4 KB page fault를 받더라도 GPU physical memory allocation은 보통 2 MB VABlock/chunk 단위로 수행한다. 따라서 CPU page evction policy와 GPU page evction policy가 일치하지 않는 문제가 발생한다.

기존 대응은 크게 prefetching, access counter 기반 migration, Zero-copy placement로 나뉜다. 그러나 prefetching은 thrashing 자체를 제거하지 못하고, access counter 방식은 static threshold에 의존하며, compiler/hardware 기반 Zero-copy 선택은 UVM의 장점인 binary transparency와 hardware portability를 약화시킨다.

== Main Idea ==
; Design Principle #1. Reduce UVM fault-handling latency through pipelined execution
: ARIADNE는 UVM fault handling의 latency bottleneck이 Copy/Eviction보다 Populate에 있다는 관찰을 사용한다. Populate, Copy, Eviction을 VABlock 간 pipeline으로 분리하여 long-latency Populate를 다른 VABlock의 Copy/Eviction과 overlap시킨다. 즉 ARIADNE는 placement policy와 fault handling pipeline을 함께 바꾸어 thrashing과 migration latency를 동시에 줄이려 한다.

; Design Principle #2. Sharing Degree: a runtime access-pattern metric leveraging thread-level information
: GPU thread execution architecture가 VABlock-level spatial locality의 runtime signal을 남긴다는 관찰을 이용하는 것이다. GPU kernel에서 thread block은 SM에 배치되고, 인접 SM들은 uTLB를 공유한다. 어떤 VABlock을 여러 uTLB/SM이 동시에 faulting한다면, 서로 다른 thread block이 그 VABlock 내부의 여러 page를 접근하고 있을 가능성이 높다. 반대로 unique uTLB가 적으면 일부 thread group만 sparse하게 접근하는 영역일 가능성이 높다. 이를 반영하는 Sharing Degree라는 runtime metric을 정의하였다. UVM driver는 kernel source나 per-thread address expression을 보지 못하지만, page fault record의 source uTLB ID는 볼 수 있다. ARIADNE는 최근 fault를 발생시킨 unique uTLB 수를 VABlock별로 추적하여, 해당 VABlock을 동시에 접근하는 SM/thread group의 수를 근사한다. 이 값이 높으면 VABlock 내부의 page utilization이 높고 dense access일 가능성이 크다는 것이 논문의 핵심 관찰이다.

; Design Principle #3. Managing CPU memory region placement between GPU memory and Zero-copy based on runtime memory access characteristics
특히 migration은 2 MB chunk 단위이고 Zero-copy는 128 B cache line 단위이므로, dense VABlock은 GPU memory에 두는 것이 유리하고 sparse VABlock은 host memory의 Zero-copy로 두는 것이 유리하다는 placement framing을 제시한다. ARIADNE는 Sharing Degree를 이용해 VABlock을 GPU memory와 Zero-copy 사이에서 동적으로 배치한다. Dense VABlock은 2 MB chunk를 GPU에 올려도 내부 활용도가 높으므로 migration이 유리하다. Sparse VABlock은 chunk를 통째로 점유하면 WCSS를 부풀리므로, evicted 이후 재접근되더라도 곧장 GPU로 다시 올리지 않고 일시적으로 Zero-copy로 둔다.

== Background ==
; VABlock과 chunk
: NVIDIA UVM은 virtual address range를 VABlock으로 관리하고, GPU physical memory는 chunk 단위로 할당한다. 논문 기준으로 보통 VABlock과 chunk는 각각 2 MB이다. Page fault는 4 KB 단위로 발생하지만 allocation과 eviction은 더 coarse-grained하게 일어난다.

; Populate, Eviction, Copy
: UVM fault handling은 GPU chunk를 준비하는 Populate, victim chunk를 host로 내보내는 Eviction, host page를 GPU memory로 복사하는 Copy로 구성된다. 논문은 Populate latency가 Copy/Eviction의 거의 두 배이며, fault batch 전체 latency를 지배한다고 분석한다.

; Zero-copy
: Zero-copy는 page를 host memory에 pin한 뒤 GPU가 remote access하도록 하는 방식이다. GPU memory bandwidth보다 느리지만, sparse VABlock을 2 MB chunk로 migrate하지 않아도 되므로 oversubscription thrashing을 줄일 수 있다.

== Design ==
; Working Chunk Set Size (WCSS) estimation
: ARIADNE는 workload가 실제로 요구하는 GPU chunk 수를 WCSS로 추적한다. 단순히 GPU-resident VABlock만 세면 방금 evict되었지만 곧 재접근될 active VABlock을 놓칠 수 있으므로, GPU resident VABlock, Zero-copy VABlock, re-access 가능성이 높은 recently evicted VABlock을 함께 포함한다. 초기 정보가 부족할 때는 GPU-wide average Sharing Degree가 낮으면 sparse workload로 보고 최근 evicted VABlock을 500 ms 동안 WCSS에 보존한다.

; Sharing Degree measurement
: 각 VABlock에 대해 최근 16개 page fault의 source uTLB ID를 circular queue에 저장하고, unique uTLB ID 개수를 Sharing Degree로 계산한다. Source uTLB ID는 UVM driver의 fault record에서 얻을 수 있으므로 hardware나 compiler modification이 필요 없다. 논문은 ATAX/GEMV 같은 sparse workload는 Sharing Degree가 주로 1이고, GEMM/HEL 같은 dense workload는 4보다 큰 값이 많으며, BFS/NW/XSB처럼 mixed pattern인 workload는 VABlock별로 다른 값을 보인다고 제시하였다. 이 Sharing Degree는 SM들에 접근하는 메모리의 spatial locality에 대한 정보를 효과적으로 보여준다고 주장한다. 따라서 '''Sparse memory의 값으로 sparse인지 아니면 dense한지 판단'''하는 것이 중요하다고 주장하였다.

; Pipelined VABlock fault handling
: 기본 UVM은 Populate, Eviction, Copy를 monolithic sequential routine으로 처리한다. ARIADNE는 Populate를 Copy와 분리하고, 다른 VABlock의 Populate/Copy/Eviction을 병렬로 실행한다. Populate는 GPU chunk allocation과 metadata 준비만 끝낸 뒤 Copy worker에 넘기고, Eviction은 free chunk가 부족해진 뒤 반응적으로 실행되는 대신 dedicated thread에서 proactive하게 실행된다. 이 구조는 Populate의 긴 latency를 다른 작업 뒤에 숨기는 것이 목적이다.

; Sharing Degree-aware eviction
: ARIADNE는 fault recency만 보는 eviction queue 대신 Sharing Degree를 반영한 priority key를 사용한다. 논문이 제시한 형태는 <code>last fault time + SD Weight * Sharing Degree / Nfault_history</code>이며, 구현에서는 fault history 길이 16과 SD Weight 100 us를 사용한다. Sharing Degree가 높은 VABlock은 GPU memory에 남길 가치가 높다고 보고 eviction priority를 낮춘다.

; Dynamic Zero-copy
: Memory demand가 GPU capacity를 넘으면 ARIADNE는 evicted 뒤 재접근된 VABlock을 즉시 GPU로 refetch하지 않고 100 ms 동안 Zero-copy state로 유지한다. 이 second-chance policy는 방금 쫓겨난 sparse VABlock이 다시 fault를 내며 eviction-refetch cycle을 만드는 것을 줄인다. 100 ms가 지나도록 재사용되지 않으면 Zero-copy state를 해제하고, 이후에도 계속 재사용되면 GPU memory로 promote될 수 있다.

; Dynamic prefetching
: Copy 단계에서는 기본 UVM prefetcher가 선택한 page를 복사하되, GPU memory 여유가 충분하거나 VABlock의 Sharing Degree가 threshold보다 높으면 VABlock 전체를 적극적으로 copy한다. 논문 구현의 Sharing Degree threshold는 3이다. 이 정책은 dense VABlock에 대해서는 future fault를 줄이고, sparse VABlock에 대해서는 불필요한 2 MB migration을 피하려는 tradeoff를 갖는다.

== Result ==
평가는 NVIDIA RTX A5000, AMD Ryzen 7700X, PCIe 4.0 x16, 64 GB DDR5, Linux 6.0, NVIDIA open-source kernel driver 535.86 환경에서 수행되었다. Workload는 Rodinia, Polybench, HeCBench, XSBench에서 가져온 10개 benchmark(2DC, ATAX, BICG, GEMM, GEMV, MVT, XSB, BFS, HEL, NW)이며, 각 benchmark의 memory footprint는 4 GB로 맞추었다. Oversubscription ratio는 no oversubscription, 130%, 175%, 300%를 사용한다.

Baseline UVM은 oversubscription에서 급격히 무너진다. 논문은 baseline UVM의 geomean execution time이 200% oversubscription에서 33.2배, 300%에서 60.7배까지 증가한다고 보고한다. AC(access counter-based migration)는 Zero-copy를 활용해 thrashing을 줄이지만 no oversubscription에서는 평균 1.3배, 최대 2.5배 느려지고, 210%와 300% oversubscription에서는 no oversubscription 대비 각각 4.7배와 8.9배 느려진다.

ARIADNE는 AC 대비 130%, 175%, 300% oversubscription에서 각각 1.9배, 2.3배, 4.0배 geomean speedup을 보인다. SUV 대비로는 같은 oversubscription ratio에서 각각 1.9배, 5.0배, 4.8배 speedup을 보인다. 또한 no-oversubscription 대비 runtime 증가는 130%, 175%, 300%에서 각각 1.6배, 1.8배, 2.3배로 보고되어, 논문은 near-linear degradation이라고 해석한다.

Mechanism-level evidence도 제시된다. Pipelined VABlock fault handling은 10개 benchmark에서 VABlock fault handling latency를 평균 17%, 최대 48%(BFS) 줄인다. Dynamic VABlock placement는 175% oversubscription에서 AC 대비 PCIe traffic을 평균 51% 수준으로 낮춘다. Breakdown에서는 Sharing Degree 기반 placement/eviction이 핵심 성능 요인이고, pipelining이 oversubscription 상황에서 정책 overhead를 숨기는 역할을 한다고 해석된다.

LLM inference 실험도 포함된다. Llama3.1 70B inference(input token length 2048)에서 ARIADNE는 AC 대비 Decode phase 4.2배, Prefill phase 1.6배 speedup을 보인다. Decode가 GEMV 중심이기 때문에 GEMV benchmark에서 큰 이득을 보인 결과와 일관된다고 설명한다.

Overhead는 작다고 보고된다. Additional metadata는 VABlock당 70 B 미만, GPU당 100 B 미만이며, 16 GB application 기준 약 560 KB이다. Sharing Degree/WCSS tracking과 Zero-copy process의 추가 latency는 최대 100 ns로, 단일 VABlock fault handling 약 20 us에 비해 무시 가능하다고 주장한다.

== Contribution ==
# UVM oversubscription에서 성능 저하의 원인을 Populate-dominated fault handling latency, 2 MB chunk granularity에 따른 WCSS amplification, last-fault-time eviction에 따른 thrashing으로 정리하였다.
# GPU thread execution architecture와 source uTLB ID를 이용해 VABlock-level spatial locality를 runtime에서 추정하는 Sharing Degree metric을 제안하였다.
# Sharing Degree와 WCSS를 기반으로 VABlock을 GPU memory와 Zero-copy 사이에서 동적으로 배치하는 runtime-only UVM management framework ARIADNE를 설계하였다.
# Populate, Copy, Eviction을 VABlock 간 pipeline으로 분리하여 UVM fault handling latency를 줄이는 driver-level execution structure를 구현하였다.
# NVIDIA open-source UVM driver 내부 수정 약 1600 LOC만으로 구현하고, hardware/compiler/application modification 없이 executable 또는 closed-source UVM application에 적용 가능한 설계를 보였다.
# 10개 GPU benchmark와 Llama3.1 70B inference에서 AC, SUV, baseline UVM 대비 성능 향상과 oversubscription scaling을 실험적으로 제시하였다.

== Criticisms ==
# Sharing Degree는 source uTLB ID를 thread group locality의 proxy로 사용하는 metric이다. 논문은 benchmark에서 utilization과의 상관을 보이지만, uTLB sharing topology, SM scheduling, phase behavior가 다른 GPU generation에서 threshold 3, history length 16, SD Weight 100 us, Zero-copy pin time 100 ms가 항상 좋은지는 제한적이다. 약간 fine-tuning되어 있는 값처럼 보이는 경향이 있다.
# 전체적으로 논문이 Collection of fine-tuned optimization으로 보이는 경향이 있다. 그러나 논문에서 제시되는 Sharing Degree와 같은 개념은 좋은 Design point이기 떄문에, Collection-of-optimizations으로 보는 것이 오히려 불필요하게 Critical한 Review일 수 있을 것 같다는 생각도 든다. 이런 Optimization work의 숙명일 수도....
# no-oversubscription 또는 memory pressure가 낮은 일부 case에서는 SUV가 ARIADNE보다 빠르다. 논문은 SUV의 compile-time range prefetching이 ARIADNE의 2 MB VABlock granularity prefetch보다 큰 data range를 미리 가져올 수 있기 때문이라고 설명한다. 즉 ARIADNE는 transparency를 얻는 대신 static program knowledge를 활용한 aggressive prefetch opportunity를 일부 잃는다.

== [[Conclusion]] ==
이 연구는 GPU UVM oversubscription을 단순 page fault overhead가 아니라 VABlock placement와 GPU thread-level locality를 함께 다루어야 하는 문제로 바라보게 만든다. ARIADNE는 source uTLB 기반 Sharing Degree, WCSS estimation, Sharing Degree-aware eviction, transient Zero-copy, pipelined fault handling을 결합하여 hardware/compiler/application 수정 없이 oversubscription 성능을 개선할 수 있음을 보였다.

[[분류: 시스템 논문]]
[[분류: IEEE HPCA]]
[[분류: GPU Unified Virtual Memory]]