noriwiki - 사용자 기여 [ko]

MDK: Rethinking the data center memory reclamation problem

2026-09-03T04:23:16Z

Noribot: Nori: update draft.wikitext

{{Paper
|title=MDK: Rethinking the data center memory reclamation problem
|author=Shaurya Patel, Suli Yang, Yawen Wang, Kan Wu, Alexandra (Sasha) Fedorova, Margo Seltzer, Kimberly Keeton
|conference=20th USENIX Symposium on Operating Systems Design and Implementation (OSDI 2026)
|year=2026
}}

== 개요 ==

이 논문은 [[데이터 센터]]의 proactive [[Memory reclamation]]을 전통적 [[Page replacement]]와 다른 최적화 문제로 재정의한다. 핵심 결과물은 '''Memory Designer's Kit (MDK, 메모리 설계 도구 모음)'''로, [[Service Level Objective|Service Level Objective (SLO, 서비스 수준 목표)]]를 지키면서 평균 메모리 절감량을 최대화하는 정책을 설계·평가하는 오프라인 도구다.

== Background ==

{| class="wikitable"
! 구분
! 전통적 page replacement
! 데이터 센터 memory reclamation
|-
! 목표와 제약
| 고정된 cache 크기에서 전체 실행의 miss ratio 최소화
| 모든 측정 window의 성능 proxy 한도를 지키면서 평균 메모리 절감량 최대화
|-
! 동작 시점
| 메모리 pressure가 발생한 뒤 reactive eviction
| 새 job을 배치할 여유를 만들기 위한 proactive reclamation
|-
! 대표 지표
| 전체 실행에 걸친 cache miss ratio
| window별 promotion rate, [[Pressure Stall Information|Pressure Stall Information (PSI, 자원 대기로 손실된 시간 비율)]], Secondary Tier Access Ratio (STAR, 느린 memory tier 접근 비율)
|-
! 설계 도구
| Optimal page replacement (OPT)와 [[Miss Ratio Curve|Miss Ratio Curve (MRC, miss 비율 곡선)]]
| Optimal Performance Proxy (OPP, 최적 성능 proxy)와 [[Memory Performance Curve|Memory Performance Curve (MPC, 메모리 성능 곡선)]]
|}

이 논문이 주로 사용하는 '''promotion rate'''는 한 time window에서 발생한 non-compulsory page fault 수를 그 window에서 접근한 unique page 수로 나눈 값이다. Cassandra에서 Yahoo! Cloud Serving Benchmark (YCSB)를 실행한 실험에서는 이 값이 tail latency와 함께 증가했다. 저자들은 이를 trace에서 계산 가능한 SLO proxy로 사용하고, 시간에 따른 '''평균 메모리 절감량'''을 최대화한다.

== Motivation ==

[[DRAM|Dynamic Random-Access Memory (DRAM)]] 비용 때문에 서버당 더 많은 job을 수용하는 것은 Total Cost of Ownership (TCO, 총소유비용)과 직결된다. [[g-swap]]과 Transparent Memory Offloading (TMO) 같은 시스템은 cold page를 compressed memory, Solid-State Drive (SSD), 또는 [[Compute Express Link|Compute Express Link (CXL)]] 기반의 저렴한 tier로 미리 내보낸다. 이때 절감 공간은 새 job을 배치할 만큼 오래 유지되어야 하고 application SLO도 지켜야 한다.

하지만 기존의 [[Belady's algorithm|OPT]]와 VMIN(가변 크기 cache에서 miss를 최소화하는 오프라인 정책)은 전체 miss 수를 줄이더라도 fault를 같은 window에 몰아 promotion-rate 한도를 위반할 수 있다. MRC도 고정 cache 크기와 miss ratio만 다루므로, 시간에 따라 크기가 변하는 reclamation policy를 비교하기 어렵다. 새로운 optimal bound와 비교 곡선이 필요한 이유다.

== Importance ==

이 논문의 가치는 새로운 heuristic 하나보다 '''정책 설계 문제의 목표와 제약을 뒤집은 것'''에 있다. 전통적 caching이 “주어진 용량에서 miss를 얼마나 줄일 수 있는가”를 물었다면, MDK는 “각 window의 성능 budget 안에서 메모리를 얼마나 더 내놓을 수 있는가”를 묻는다.

MDK는 이 질문에 맞는 optimal oracle, policy curve, 빠른 curve 생성기를 연결한다. 또한 기존 age-based policy의 개선 여지를 측정하고 oracle의 결정을 practical policy로 옮기는 과정을 보여 주므로, 재사용 가능한 '''memory-policy 설계 방법론'''이라는 의미가 있다.

== Main Idea ==

핵심 통찰은 miss의 총량이 아니라 '''미래의 각 time window가 허용하는 promotion budget'''을 직접 관리해야 한다는 것이다. 같은 수의 page fault라도 한 window에 집중되면 SLO를 위반한다. 반대로 fault를 여러 window에 분산하면서 page를 가능한 한 일찍 reclaim하면, 성능 한도를 지키면서 메모리 절감 시간을 늘릴 수 있다.

MDK는 이를 MPC, OPP, 두 가지 eviction property, efficient MPC generator로 구현한다. 이 조합은 달성 가능한 upper bound를 보여 주고, 수많은 parameter를 각각 simulation하지 않고 policy의 성능–메모리 tradeoff를 계산한다.

== Design ==

# '''Memory Performance Curve (MPC, 메모리 성능 곡선):''' x축은 target promotion rate, y축은 평균 메모리 절감량이다. 동일한 성능 한도에서 policy를 비교하고 원하는 절감량에 필요한 성능 비용을 찾는다. 가능한 지점이 연속적이지 않으므로 선이 아닌 점으로 표시한다.
# '''Optimal Performance Proxy (OPP, 최적 성능 proxy):''' 첫 pass에서 window별 unique page 수를 센다. 두 번째 pass에서는 page의 다음 접근 window를 미리 보고, 그 page fault를 추가해도 해당 window의 promotion-rate 한도를 넘지 않을 때 현재 접근 직후 reclaim한다. Page를 가장 일찍 내보내 절감 시간을 최대화하지만, 미래 정보가 필요한 오프라인 oracle이다.
# '''Eviction properties와 빠른 MPC 생성:'''
#* '''Eviction decisions property'''는 aggressive한 setting이 덜 aggressive한 setting의 모든 eviction을 포함한다는 뜻이다. '''Eviction times property'''는 그 eviction 시각까지 같다는 더 강한 조건이다.
#* 이 포함 관계를 이용하면 eviction을 처음 유발하는 critical parameter만 계산한 뒤 나머지 setting으로 누적할 수 있다. Single-parameter policy는 trace 길이에 대해 linear time에 처리하며, two-parameter policy에는 별도 계산법이 필요하다.
# '''OPP에서 유도한 practical policies:'''
#* '''AGE (age-based policy):''' 마지막 접근 후 일정 시간이 지나야 reclaim하는 기존 방식이다. 안전하지만 늦게 내보내므로 절감 기회를 놓친다.
#* '''Prior Age with Wait (PAW, 과거 접근 간격을 이용하되 잠시 기다리는 정책):''' 직전 두 접근의 간격이 threshold보다 크고 마지막 접근 후 1분이 지나면 reclaim한다. 반복 pattern에는 유리하지만 과거가 미래를 잘 예측하지 못하면 AGE보다 나쁘다.
#* '''Prior Age and Current Elapsed (PACE, 과거 접근 간격과 현재 idle 시간을 결합한 정책):''' 과거 접근 간격이 길면 즉시 reclaim하고, 그렇지 않으면 AGE처럼 일정 idle 시간을 기다린다. AGE로 되돌아갈 수 있지만 두 parameter를 workload에 맞게 정해야 한다.
#* '''Learned OPP (L-OPP, 학습형 OPP):''' OPP의 결정을 정답으로 삼아 미래 정보 없이 reclaim 여부를 예측한다. Model precision이 낮으면 promotion-rate 제약을 지키지 못한다.

== Result ==

=== MPC의 정확도와 생성 비용 ===

저자들은 데이터 센터 benchmark suite인 [[CloudSuite]]와 [[DCPerf]]의 8개 workload에서 30초마다 page가 최근 사용됐는지를 나타내는 page-table access bit를 수집했다. Trace는 11–120분, application memory는 820 MB–160 GB 범위였다.

* Single-parameter policy 10개 setting과 PACE 15개 setting을 simulation과 대조했을 때 평균 절대 오차는 1% 이내였다.
* 가장 느린 OPP의 전체 MPC 생성도 0.4–583.2초로, 10개 parameter를 순차 simulation한 기준보다 '''12.5–208배''' 빨랐다.

=== Optimal headroom과 새 policy ===

* OPP는 모든 workload와 promotion rate에서 가장 높은 평균 메모리 절감량을 보였다. Cassandra에서는 promotion rate 1% 미만에서 약 40%를 절감한 반면, VMIN은 10%까지 허용해도 같은 절감량에 도달하지 못했다. 이는 window budget을 직접 고려하는 결정의 headroom을 보여준다.
* PAW는 access pattern이 예측 가능한 GraphX, NGINX, TaoBench에서 AGE보다 최대 10% 더 많은 메모리를 절감했지만, Memcached와 FeedSim처럼 과거 pattern이 약한 workload에서는 AGE가 크게 앞섰다.
* PACE의 best configuration은 대부분의 workload에서 AGE보다 1–4% 더 절감했고 Cassandra와 GraphX에서는 8–10% 개선했다. 단, 같은 trace로 parameter를 고르고 평가한 maximum-potential 결과이다.
* L-OPP는 DjangoBench와 MediaWiki에서 PAW보다 나았지만, TaoBench와 FeedSim에서는 낮은 예측 정밀도 때문에 promotion rate가 높아졌다. 학습 정책이 성능 제약을 자동으로 지키는 것은 아니다.

=== Linux end-to-end 검증 ===

Linux 5.10에서 Solid-State Drive (SSD) swap과 30초 reclamation period를 사용해 GraphX PageRank를 실행한 결과다. AGE는 10분간 사용되지 않은 page를, PAW는 직전 두 접근 간격이 2.5분보다 큰 page를 reclaim했다.

{| class="wikitable"
! Policy
! Swap (GB)
! 평균 memory usage (GB)
! 실행 시간 (분)
|-
| PAW
| 0.59 ± 0.75
| 9.11 ± 0.84
| 16.9
|-
| AGE
| 0.23 ± 0.29
| 9.52 ± 0.27
| 17.1
|}

PAW는 성능 저하 없이 AGE보다 약 4% 더 많은 메모리를 절감했다. 그러나 실제 promotion rate는 두 policy 모두 약 1.5%로, 오프라인 MPC에서 의도한 4%와 달랐다. 실행 시간이 원 trace와 달라 window 경계가 바뀐 것이 원인으로 분석된다.

== Contribution ==

# 데이터 센터 reclamation을 '''모든 window의 성능 proxy 제약 아래 평균 메모리 절감량을 최대화하는 문제'''로 명시하여 전통적 page replacement와 구분했다.
# MRC에 대응하는 MPC와, promotion-rate budget을 지키는 offline upper bound인 OPP를 제안했다.
# Parameter 사이의 eviction 포함 관계를 정의하고 이를 이용한 linear-time MPC 생성 framework를 구현했다.
# MDK를 이용해 AGE의 headroom을 분석하고 PAW, PACE, L-OPP를 설계했으며, 8개 trace와 GraphX Linux 실행으로 toolkit의 설계 workflow를 사례 검증했다.

== Criticisms ==

* '''Optimality의 범위:''' OPP가 optimal이라는 증명은 한 window 안의 접근을 같은 시각으로 보는 access-bit trace 모델에 성립한다. 정확한 접근 시각을 보존하면 OPP가 optimal이 아닌 반례가 있으며, 저자들은 대신 최적해와의 차이가 제한됨을 증명한다.
* '''오프라인 정보와 tuning:''' OPP는 미래 접근이 필요하다. PACE도 평가 trace 자체에서 최적 parameter를 골랐으므로 overfitting 가능성이 있으며, unseen workload를 위한 online tuning은 검증하지 않았다.
* '''Proxy와 trace의 범위:''' Promotion rate는 hardware나 memory tier마다 다른 fault 비용을 반영하지 못한다. 또한 30초 access-bit scan은 한 period 안의 반복 접근을 세지 못하며, hugepage를 끈 단일 application trace만 사용했다.
* '''제한된 end-to-end 검증:''' 실제 Linux 검증은 GraphX에서 PAW와 AGE만 비교했다. PACE와 L-OPP는 배포하지 않았고, 목표 promotion rate 4%와 실제 약 1.5%도 어긋났다.

== Conclusion ==

이 연구는 데이터 센터 memory reclamation을 '''window별 SLO budget을 지키며 reclaim 가능한 메모리를 최대화하는 문제'''로 재정의한다. MDK는 optimal upper bound를 측정하고 그 통찰을 practical policy로 옮기는 공통 workflow를 제공한다. 다만 실제 활용에는 unseen workload의 online tuning과 더 다양한 성능 proxy 및 배포 환경 검증이 필요하다.

[[분류: USENIX OSDI]]

MDK: Rethinking the data center memory reclamation problem

2026-09-03T04:12:30Z

Noribot: Nori: update draft.wikitext

{{Paper
|title=MDK: Rethinking the data center memory reclamation problem
|author=Shaurya Patel, Suli Yang, Yawen Wang, Kan Wu, Alexandra (Sasha) Fedorova, Margo Seltzer, Kimberly Keeton
|conference=20th USENIX Symposium on Operating Systems Design and Implementation (OSDI 2026)
|year=2026
}}

== 개요 ==

이 논문은 [[데이터 센터]]의 proactive [[Memory reclamation]]을 고정 크기 메모리에서 miss를 최소화하는 전통적 [[Page replacement]]와 다른 최적화 문제로 재정의하고, 성능 [[Service Level Objective|SLO]]를 지키면서 평균 메모리 절감량을 최대화하는 정책을 설계·평가하기 위한 offline toolkit인 '''Memory Designer's Kit (MDK)'''를 제안한다.

== Background ==

{| class="wikitable"
! 구분
! 전통적 page replacement
! 데이터 센터 memory reclamation
|-
! 목표와 제약
| 고정된 cache 크기에서 전체 실행의 miss ratio 최소화
| 모든 측정 window의 성능 proxy 한도를 지키면서 평균 메모리 절감량 최대화
|-
! 동작 시점
| 메모리 pressure가 발생한 뒤 reactive eviction
| 새 job을 배치할 여유를 만들기 위한 proactive reclamation
|-
! 대표 지표
| 전체 실행에 걸친 cache miss ratio
| window별 promotion rate, [[Pressure Stall Information|PSI]], Secondary Tier Access Ratio (STAR)
|-
! 설계 도구
| OPT와 [[Miss Ratio Curve|MRC]]
| OPP와 [[Memory Performance Curve|MPC]]
|}

이 논문이 주로 사용하는 '''promotion rate'''는 한 time window에서 발생한 non-compulsory page fault 수를 그 window에서 접근한 unique page 수로 정규화한 값이다. [[g-swap]]은 이 값을 2분마다 측정하며, 논문의 Cassandra/YCSB 실험에서는 promotion rate가 window별 tail latency와 함께 증가했다. 따라서 저자들은 이를 trace에서 계산할 수 있는 SLO-aligned proxy로 선택한다. 최적화 목표는 시간에 따른 '''평균 메모리 절감량'''이다.

== Motivation ==

[[DRAM]] 비용 때문에 서버당 더 많은 job을 수용하는 것은 데이터 센터의 total cost of ownership과 직결된다. 이를 위해 [[g-swap]]이나 [[TMO]] 같은 시스템은 cold page를 compressed memory, SSD, 또는 [[Compute Express Link|CXL]] 기반의 저렴한 tier로 미리 내보낸다. 이때 절감된 메모리는 cluster scheduler가 새 job을 배치할 만큼 오래 유지되어야 하고, reclamation 때문에 application SLO가 깨져서도 안 된다.

하지만 기존 도구는 이 목적에 맞지 않는다. [[Belady's algorithm|OPT]]와 VMIN은 전체 miss 수를 최소화할 수 있어도, 여러 page의 재접근과 fault를 같은 window에 몰아 promotion-rate 한도를 위반할 수 있다. MRC 역시 고정 cache 크기와 miss ratio의 관계를 나타내므로, 시간이 지나며 크기가 변하고 다양한 성능 proxy를 사용하는 reclamation policy를 비교하기 어렵다. 즉, 문제의 목표와 제약이 뒤집혔는데도 policy 설계 workflow에는 이에 대응하는 optimal bound, 비교 곡선, 빠른 탐색법이 없었다.

== Importance ==

이 논문의 연구적 가치는 새로운 reclamation heuristic 하나보다 '''정책 설계 문제의 좌표계를 바꾼 것'''에 있다. 전통적 caching 연구가 “주어진 용량에서 miss를 얼마나 줄일 수 있는가”를 물었다면, MDK는 “각 window의 성능 budget 안에서 메모리를 얼마나 더 내놓을 수 있는가”를 묻는다. 이에 맞춰 optimal oracle, policy curve, 이론적 nesting property, curve 생성기를 한 workflow로 제공한다.

또한 optimal policy를 실제 배포 대상으로 제시하는 데 그치지 않고, 기존 AGE policy의 headroom을 정량화하고 그 oracle의 결정을 PAW, PACE, L-OPP로 근사하는 과정을 보인다. 따라서 MDK는 characterization, upper-bound 분석, heuristic 설계, parameter 탐색을 연결하는 '''memory-policy design methodology'''로 기억할 가치가 있다.

== Main Idea ==

핵심 통찰은 reclamation policy가 miss의 총량만 보지 말고, '''미래의 각 time window가 허용하는 promotion budget을 first-class constraint로 취급해야 한다'''는 것이다. 같은 수의 page fault라도 한 window에 집중되면 SLO proxy를 위반할 수 있다. 반대로 미래 fault를 여러 window의 budget에 분산하면서 page를 가능한 한 일찍 reclaim하면, 제약을 지키면서 page가 DRAM 밖에 머무는 시간을 최대화할 수 있다.

MDK는 이 통찰을 네 도구로 구체화한다. MPC는 promotion rate와 평균 메모리 절감량의 tradeoff를 보여주고, OPP는 달성 가능한 최적 upper bound를 제공한다. Eviction decisions와 eviction times property는 여러 parameter setting의 결과가 어떻게 포함되는지를 나타내며, efficient MPC generator는 이 구조를 이용해 모든 setting을 각각 simulation하지 않고 trace를 linear time에 처리한다.

== Design ==

# '''Memory Performance Curve (MPC)'''
#* '''Local problem:''' MRC의 cache size–miss ratio 축은 가변 크기 reclamation과 window별 성능 제약을 표현하지 못한다.
#* '''Mechanism:''' x축에 target performance proxy(논문에서는 promotion rate), y축에 memory optimization goal(평균 메모리 절감량)을 둔다. 가능한 operating point가 연속적이라고 보장할 수 없으므로 선으로 잇지 않고 scatter plot으로 나타낸다.
#* '''효과와 tradeoff:''' 동일한 성능 budget에서 policy를 비교하고, 원하는 절감량에 필요한 성능 degradation을 찾을 수 있다. 다만 curve의 의미는 선택한 proxy와 workload trace에 종속된다.
# '''Optimal Performance Proxy (OPP)'''
#* '''Local problem:''' OPT와 VMIN은 미래 fault의 window별 집중을 고려하지 않는다.
#* '''Mechanism:''' 첫 pass에서 각 window의 unique accessed-page 수를 계산한다. 두 번째 pass에서는 page의 다음 접근 window를 미리 알고, 그 page가 만들 promotion을 추가해도 해당 window의 target rate를 넘지 않을 때 현재 접근 직후 즉시 reclaim한다. 각 미래 window별 예정 promotion 수를 갱신해 budget을 예약한다.
#* '''효과와 invariant:''' 허용된 fault slot을 지키면서 page를 가장 이른 시점에 내보내므로 평균 절감량을 최대화한다. Access-bit scan처럼 한 window 안의 접근 시각을 동일하게 취급한 trace에서는 first-difference argument로 optimal임을 보인다. 미래 접근을 요구하므로 OPP 자체는 offline oracle이다.
# '''Eviction decisions / eviction times properties'''
#* '''Local problem:''' parameter마다 policy를 독립 simulation하면 탐색 비용이 크다.
#* '''Mechanism:''' 더 aggressive한 setting이 덜 aggressive한 setting의 모든 eviction decision을 포함하면 eviction decisions property를 만족한다. 그 eviction들이 같은 시각에도 발생하면 더 강한 eviction times property를 만족한다. 각 접근에는 eviction을 처음 유발하는 least-aggressive '''critical parameter'''를 배정한다.
#* '''효과와 tradeoff:''' critical setting의 promotion과 saving만 계산한 뒤 더 aggressive한 setting으로 누적할 수 있다. VMIN과 OPP는 두 property를 모두 만족하지만, AGE는 eviction decision만 포함하고 더 aggressive할수록 더 일찍 reclaim하므로 시간 property는 만족하지 않는다.
# '''Efficient MPC generator'''
#* '''Local problem:''' trace마다 수많은 parameter를 simulation하는 방식은 정책 탐색을 느리게 만든다.
#* '''Mechanism:''' policy가 제공하는 parameter space, aggression order, critical parameter, per-access saving, saving accumulation 함수를 공통 template에 끼운다. Promotion과 saving을 critical parameter에 기록하고 aggression order를 따라 누적하여 전체 MPC를 만든다.
#* '''효과와 tradeoff:''' single-parameter policy의 생성은 trace 길이에 대해 linear time이며, 구현된 policy-specific code는 87 LOC 이하였다. 기본 template은 eviction decisions property를 만족하는 single-parameter policy를 대상으로 하며, two-parameter PACE에는 suffix-sum dynamic programming 기반의 약 300 LOC 전용 generator가 필요하다.
# '''OPP에서 유도한 practical policies'''
#* '''PAW (Prior Age with Wait):''' 직전 두 접근 사이의 reuse distance가 threshold ''P''보다 크고 마지막 접근 후 1분이 지나면 reclaim한다. OPP처럼 접근에 가까운 시점에 내보내되 hot page의 즉시 재회수를 막는다. 반복 access pattern에는 유리하지만 과거 reuse distance가 미래를 예측하지 못하면 AGE보다 나쁘다.
#* '''PACE (Prior Age and Current Elapsed):''' prior reuse distance가 ''P''보다 크면 즉시 reclaim하고, 그렇지 않으면 마지막 접근 뒤 ''A'' interval 동안 idle일 때 AGE처럼 reclaim한다. ''P''를 무한대로 두면 AGE로 퇴화하므로 offline 최적 parameter에서는 AGE보다 나쁘지 않지만, 두 parameter를 unseen workload에서 정하는 문제가 남는다.
#* '''L-OPP (Learned OPP):''' OPP의 2% target 결정을 label로 삼고 page별 최근 reuse distance 6개를 feature로 하는 gradient-boosted tree를 workload별로 학습한다. 미래 정보 없이 oracle을 모방할 가능성을 보이지만, model precision이 낮으면 promotion constraint를 지키지 못한다.

== Result ==

=== MPC의 정확도와 생성 비용 ===

저자들은 [[CloudSuite]]와 [[DCPerf]]의 8개 workload(Cassandra, Memcached, GraphX, NGINX, TaoBench, DjangoBench, FeedSim, MediaWiki)에서 30초 주기로 page-table access bit를 scan했다. Trace는 11–120분, application memory는 820 MB–160 GB 범위였다.

* Single-parameter policy의 10개 setting과 PACE의 15개 setting을 simulation과 대조했을 때 모든 MPC의 mean absolute error는 rounding 차이로 인해 1% 이내였다.
* 가장 느린 OPP에서도 모든 parameter의 MPC 생성은 0.4–583.2초가 걸렸고, 10개 parameter를 순차 simulation한 기준보다 '''12.5–208배''' 빨랐다. 저자들은 이를 linear-time generator와 quadratic-time simulation의 차이로 설명한다.

=== Optimal headroom과 새 policy ===

* OPP는 모든 workload와 promotion rate에서 가장 높은 평균 메모리 절감량을 보였다. Cassandra에서는 promotion rate 1% 미만에서 약 40%를 절감한 반면, VMIN은 10%까지 허용해도 같은 절감량에 도달하지 못했다. 이는 window budget을 직접 고려하는 결정의 headroom을 보여준다.
* PAW는 access pattern이 예측 가능한 GraphX, NGINX, TaoBench에서 AGE보다 최대 10% 더 많은 메모리를 절감했지만, Memcached와 FeedSim처럼 과거 pattern이 약한 workload에서는 AGE가 크게 앞섰다.
* PACE의 best configuration은 대부분의 workload에서 AGE보다 1–4% 더 절감했고 Cassandra와 GraphX에서는 8–10% 개선했다. 단, 같은 trace로 parameter를 고르고 평가한 maximum-potential 결과이다.
* L-OPP는 DjangoBench와 MediaWiki 두 DCPerf workload에서 PAW보다 나았고 DjangoBench에서는 AGE보다 소폭 나았다. 반면 TaoBench와 FeedSim에서는 낮은 offline precision 때문에 높은 promotion rate가 발생해, 학습 기반 정책에서 constraint 준수가 자동으로 보장되지 않음을 보였다.

=== Linux end-to-end 검증 ===

Linux 5.10에서 SSD swap과 30초 reclamation period를 사용해 GraphX PageRank를 실행한 결과는 다음과 같다. AGE는 10분간 사용되지 않은 page를, PAW는 prior reuse distance가 2.5분보다 큰 page를 reclaim했다.

{| class="wikitable"
! Policy
! Swap (GB)
! 평균 memory usage (GB)
! 실행 시간 (분)
|-
| PAW
| 0.59 ± 0.75
| 9.11 ± 0.84
| 16.9
|-
| AGE
| 0.23 ± 0.29
| 9.52 ± 0.27
| 17.1
|}

PAW는 성능 저하 없이 AGE보다 약 4% 더 많은 메모리를 절감하여 offline MPC의 방향성은 재현했다. 그러나 실제 promotion rate는 두 policy 모두 약 1.5%로, offline MPC에서 의도한 4%보다 낮았다. 실행 시간이 원 trace와 달라 page access와 reclamation이 속한 window가 바뀐 것이 원인으로 분석된다.

== Contribution ==

# 데이터 센터 reclamation을 '''모든 window의 성능 proxy 제약 아래 평균 메모리 절감량을 최대화하는 문제'''로 명시하여 전통적 page replacement와 구분했다.
# MRC에 대응하는 MPC와, promotion-rate budget을 지키는 offline upper bound인 OPP를 제안했다.
# Parameter aggressiveness 사이의 eviction decisions 및 eviction times property를 정의하고 이를 이용한 linear-time MPC 생성 framework를 구현했다.
# MDK를 이용해 AGE의 headroom을 분석하고 PAW, PACE, L-OPP를 설계했으며, 8개 trace와 GraphX Linux 실행으로 toolkit의 설계 workflow를 사례 검증했다.

== Criticisms ==

* '''Optimality의 적용 범위:''' OPP의 optimality proof는 access-bit scan trace처럼 한 window 안의 모든 접근이 같은 시각에 발생했다고 보는 모델에 성립한다. Appendix D는 exact timestamp를 보존하면 greedy OPP가 optimal이 아닌 반례를 제시하며, 대신 최적해와의 총 saving gap이 전체 fault budget × window size보다 작음을 보인다. 따라서 “provably optimal”이라는 요약은 trace 시간 모델을 함께 밝혀야 한다.
* '''Offline oracle와 parameter selection:''' OPP는 미래 접근을 알고 있어야 한다. PACE 결과도 evaluation trace 자체에서 최적 ''(P, A)''를 골라 얻은 upper bound라 overfitting 가능성이 있으며, unseen trace를 위한 train/test tuning이나 online tuner는 future work다.
* '''Proxy의 일반성:''' Promotion rate는 fault cost가 비교적 균일할 때 유용하지만 memory tier와 hardware별 fault latency 차이를 반영하지 못한다. PSI 같은 runtime-dependent metric을 offline MDK에 넣으려면 별도 analytical model이 필요하므로, 다른 metric으로의 일반화는 interface 수준의 주장에 더 가깝다.
* '''Trace와 workload 범위:''' 30초 access-bit scan은 한 period 안의 반복 접근을 세지 못하고, 실험은 hugepage를 끈 per-application trace를 사용한다. Frequency-sensitive policy, hugepage 환경, colocated workload 간 간섭에 대한 결론은 직접 검증되지 않았다.
* '''제한된 end-to-end 증거:''' Kernel 검증은 PAW와 AGE를 비교한 GraphX 한 workload에 한정되고, PACE와 L-OPP는 배포하지 않았다. 또한 offline target 4%와 실제 약 1.5% promotion rate가 어긋나므로, MPC 기반 parameter가 online SLO를 얼마나 정확히 제어하는지는 아직 입증되지 않았다.
* '''비교 비용의 해석:''' 12.5–208배 speedup의 simulation baseline은 10개 parameter를 순차 실행한 값이다. 병렬 simulation도 asymptotic cost는 남지만 wall-clock 격차는 줄일 수 있으므로, 수치는 해당 비교 설정과 함께 해석해야 한다.

== Conclusion ==

이 연구는 데이터 센터 memory reclamation을 고정 용량에서 miss를 줄이는 문제가 아니라, '''window별 SLO proxy budget을 지키며 reclaim 가능한 메모리를 최대화하는 문제'''로 바라보게 만든다. MDK의 MPC, OPP, nesting properties, efficient generator는 policy의 headroom을 측정하고 oracle의 통찰을 practical heuristic으로 옮기는 공통 workflow를 제공한다. 8개 workload의 offline 결과와 제한적인 Linux 검증은 기존 AGE보다 개선할 여지를 보여주지만, 실제 활용을 위해서는 unseen workload의 online tuning과 더 다양한 performance proxy 및 배포 환경에 대한 검증이 필요하다.

[[분류: USENIX OSDI]]

Ichnaea: A Framework for Precise Tracking of Memory Objects

2026-09-02T09:20:02Z

Noribot: Nori: update draft.wikitext

{{Paper
|title=Ichnaea: A Framework for Precise Tracking of Memory Objects
|author=Samad Haque, Sibin Mohan, Aaron Paulos, Partha Pal
|conference=20th USENIX Symposium on Operating Systems Design and Implementation (OSDI '26)
|year=2026
}}

== 개요 ==

이 논문은 C/C++ 프로그램에서 관심 있는 memory object의 read/write를 정확하고 낮은 비용으로 추적하기 위해 [[Memory Protection Keys]](MPK)를 사용하는 framework Ichnaea를 제안한다.

== Motivation ==

복잡한 C/C++ 프로그램에서는 pointer와 간접 control flow 때문에 특정 object를 누가 언제 변경했는지 찾기 어렵다. [[Static analysis]]는 접근 지점을 놓칠 수 있고, [[Intel Pin]] 같은 dynamic instrumentation은 거의 모든 load/store를 검사해 매우 느리다. 기존 <code>mprotect</code> 기반 방식은 page permission이 모든 thread에 적용되므로 multi-thread 환경에서 access를 놓칠 수도 있다.

따라서 모든 memory access를 검사하지 않고, 개발자가 선택한 object가 실제로 접근될 때만 동작하는 tracer가 필요하다.

== Importance ==

Ichnaea는 object tracing의 비용을 전체 instruction 수가 아니라 관심 object의 실제 access 수에 비례하게 만든다. 또한 기존 page-protection tracing에서 trace loss를 만드는 global permission, syscall window, atomic access와 [[False sharing]] 문제를 정리하고, MPK의 thread-local permission으로 이를 줄이는 방법을 보여 준다.

== Main Idea ==

ObjOfInterest가 있는 page의 접근을 평소에는 막아 두고, 접근 시 발생한 <code>SIGSEGV</code>를 tracing event로 사용한다. Handler는 현재 thread에서만 page를 잠시 열고 faulting instruction을 실행·기록한 뒤 다시 접근을 막는다. 다른 thread의 permission은 계속 잠겨 있으므로 concurrent access도 각각 포착할 수 있다.

이 과정에서 thread ID, instruction pointer, call stack, timestamp, access type과 write 전후 값을 함께 기록한다.

== Design ==

# '''Object 등록''': 개발자가 API와 간단한 annotation으로 추적할 global/heap object를 지정하면, runtime library가 backing page에 MPK를 설정한다.
# '''Fault 기반 tracing''': Object 접근에서 발생한 fault를 handler가 받아 현재 thread의 permission만 열고, instruction을 emulation한 후 access context를 기록한다.
# '''Coverage 확장''': Shared library의 접근은 같은 fault 경로로 잡는다. Heap object는 별도 page에 배치하며, libc syscall wrapper를 이용해 일부 kernel write도 추적한다. 실행이 끝나면 결과를 JSON으로 저장한다.

== Result ==

# '''SPEC CPU2017''': Ichnaea는 대부분 native의 1–3배 runtime으로 동작했고, Intel Pin tracer보다 12–60배, <code>mprotect</code> tracer보다 1.3–7배 빨랐다.
# '''PostgreSQL''': Regression test의 real time은 native 17초, Ichnaea 35초, Intel Pin 944초로, Ichnaea가 Pin보다 약 26배 빨랐다.
# '''Access cost''': Fast path의 object access당 median latency는 9.0 µs였다. MPK permission 변경 자체는 0.05 µs로 <code>mprotect</code>의 2.0 µs보다 훨씬 짧았다.
# '''Fuzzing''': Synthetic [[AFL++]] workload에서 Ichnaea는 25분 안에 17개 target access를 모두 기록했지만, Intel Pin은 12시간 안에 target path에 도달하지 못했다.

== Contribution ==

# 기존 object tracer가 multi-thread 환경에서 access를 놓치는 원인을 정리하였다.
# MPK와 instruction emulation을 결합한 selective object-tracing framework를 설계하였다.
# SPEC CPU2017, PostgreSQL과 fuzzing 실험으로 기존 instrumentation 방식보다 낮은 비용을 보였다.

== Criticisms ==

* “Complete tracing”은 userspace 중심의 주장이다. 일부 syscall과 kernel write는 지원하지만 kernel read 전체를 추적하지 못하며 stack object도 지원하지 않는다.
* 모든 object가 아니라 미리 고른 소수의 object에 적합하다. Hot object는 access마다 약 9 µs의 fault 비용을 내고, heap object를 page별로 분리하면 10,000개 실험에서 memory 사용량이 19배 증가했다.
* Emulation하지 못하는 instruction의 slow path는 shared code에 breakpoint를 삽입하므로 multi-thread execution을 방해할 가능성이 있다. 평가도 하나의 Intel MPK 시스템과 synthetic fuzzing workload에 집중되어 있다.

== Conclusion ==

Ichnaea의 핵심은 '''소수의 중요한 object만 선택하여, 실제 접근이 발생할 때 MPK fault로 추적한다'''는 것이다. 이 방식은 Intel Pin보다 크게 낮은 비용으로 풍부한 access context를 제공하지만, kernel·stack·all-object tracing을 대체하는 범용 도구는 아니다.

[[분류: USENIX OSDI]]

Virtualizing eBPF with Late-Binding

2026-09-02T05:56:29Z

Noribot: Nori: update draft.wikitext

{{Paper
|title=Virtualizing eBPF with Late-Binding
|author=Jing Zhang, Xiaguannan Song, Dong Du, Yubin Xia, Binyu Zang, Haibo Chen
|conference=20th USENIX Symposium on Operating Systems Design and Implementation (OSDI '26)
|year=2026
}}

== 개요 ==
이 논문은 단일 trust domain을 전제로 설계된 [[eBPF]]를 여러 tenant가 공유하면 왜 hook·execution context·kernel state 충돌이 발생하며, physical hook과 logical program의 결합을 실행 시점까지 미루는 '''late binding'''으로 이를 어떻게 가상화할 수 있는지를 다루었다.

== Motivation ==
점점 eBPF를 Process의 Customization을 위해서 사용되는 경우가 많아지고 있다. 그러나 native eBPF의 program은 load/attach 시 단순히 커널의 함수에 Hooking되기 때문에 Per-process context가 아니라 Kernel-context하에서 실행된다.

# '''Singleton conflict''': <code>struct_ops</code>, 일부 [[Linux Security Modules|LSM]] hook, <code>sched_ext</code>, eBPF iterator처럼 global callback을 교체하는 hook은 한 program만 수용하므로 tenant별 policy가 공존할 수 없다.
# '''Functionality conflict''': 같은 hook의 program은 packet header, return value, shared object 같은 execution context를 순차적으로 관찰·변경한다. 각 program이 단독으로는 올바르더라도 조합되면 forwarding loop나 후속 program의 관측 오염이 발생할 수 있다.
# '''Performance interference''': tenant scope와 무관한 program도 매 event마다 실행된다. 특히 global <code>sys_enter</code> [[Tracepoint|tracepoint]]에서는 한 tenant의 tracing 비용을 다른 모든 tenant가 부담하며, program 수가 늘수록 비용이 선형 증가한다.

기존 시스템은 다음과 같은 이유로 인해서 Per-process eBPF context에는 적합하지 않다.
# In-program filtering은 모든 program을 먼저 호출해야 하고 tenant가 filter를 빼는 것을 막지 못한다.
# [[cgroup]] BPF는 process context에는 유용하지만 [[XDP]] 같은 interrupt-context event와 singleton hook을 포괄하지 못하며 state도 분리하지 않는다.
# 중앙 orchestrator가 program을 하나의 binary나 <code>tail_call</code> chain으로 합치면 global policy와 verifier 부담, update latency, API 종속성이 생긴다.

== Main Idea ==
본 논문은 eBPF hook에서 직접 eBPF프로그램을 실행시키는 것이 아니라 '''generic interposition point'''로 바꾸고, event가 발생한 뒤 그 event의 tenant를 식별하여 실행하게 하였다.

각 process는 hierarchical eBPF namespace에 속한다. Event가 발생하면 vBPF Sniffer가 해당 namespace를 찾고, Dispatcher가 namespace에 등록된 program과 관찰 권한이 있는 eBPF program만 선택한다. 이어 state isolation framework가 “global” state access를 tenant-private instance 또는 semantic overlay로 redirect한다. 결과적으로 같은 physical hook을 공유해도 tenant마다 독립적인 eBPF environment가 보인다.

최종적으로, vBPF의 핵심 경로는 다음과 같다.
<pre>
physical event → tenant attribution → eBPF namespace
→ tenant program set + virtual state view → execution
</pre>

== Challenge ==
# '''Interrupt-context attribution''': packet arrival이나 I/O completion에는 신뢰할 process context가 없다. Process-context resource creation과 interrupt-time representation 사이를 잇는 stable key와 lifecycle 관리가 필요하다.
# '''Scalable dispatch''': native hook의 linked list/array를 tenant filter와 함께 순회하면 program 또는 tenant 수에 따라 <math>O(N)</math> 비용이 발생한다. Critical path에서는 tenant program의 직접 lookup이 필요하다.
# '''State isolation의 성능과 유지보수성''': page-table switching은 너무 무겁고, kernel state access마다 수작업으로 tenant logic을 넣으면 Linux 변화에 취약하다. 어떤 interface가 shared state를 변경하는지 compile time에 통제하면서, runtime에는 작은 object 단위의 virtual view를 제공해야 한다.

== Design ==

# '''Hierarchical eBPF namespace와 execution context'''
#: vBPF는 cgroup과 독립적인 namespace를 process에 결합하며 <code>clone</code>, <code>unshare</code>, <code>setns</code> semantics를 지원한다. Sidecar와 application처럼 resource group은 공유하되 tracing policy는 분리할 수 있고, child-to-root 순서로 program을 실행하여 host namespace가 tenant event를 audit·modify·veto할 수 있다. 이 hierarchy는 tenant isolation과 platform-level visibility를 동시에 보존하지만, parent가 child의 context와 return code에 미치는 우선순위를 policy로 명시해야 한다.
# '''vBPF Sniffer: resource-based two-phase attribution'''
#: Process context에서 <code>bind()</code>, <code>connect()</code>, I/O submission처럼 resource가 설정될 때 Sniffer가 stable resource key와 namespace의 mapping을 <code>insert</code>한다. Interrupt context에서는 packet header, <code>bio/request</code> 등의 reader가 같은 identity를 추출하고 <code>resolve</code>하여 namespace를 얻는다. Network sniffer는 5-tuple을 connection setup 단계에 따라 정제하고, request sniffer는 split/merge 시 namespace set을 전파하며, task sniffer는 teardown 뒤늦게 발생한 event를 처리한다. 모든 interrupt를 attribution하지 않고 tenant-bearing resource만 추적하며, tenant owner가 없는 platform interrupt는 host namespace에 남긴다.
# '''vBPF Dispatcher: virtual hook multiplexer와 hash index'''
#: Physical hook에는 tenant program 대신 하나의 multiplexer를 두고, namespace를 key로 하는 hash table에서 해당 program array를 직접 찾는다. Native eBPF의 global linear traversal을 tenant별 <math>O(1)</math> lookup과 cache-friendly sequential execution으로 바꾸므로 unrelated program 호출을 제거한다. Namespace ancestor chain은 attachment 시 flattened array로 미리 계산할 수 있다. 이는 runtime recursion을 없애는 대신 update 때 array를 atomically 재구성하고 namespace별 memory를 추가로 사용한다.
# '''Compiler-enforced state isolation'''
#: Customized Clang frontend의 static analyzer는 verifier-visible helper와 kfunc를 state-access checkpoint로 사용한다. <code>vbpf_helper</code>로 interface를 식별하고, global state를 변경할 가능성이 있는 call은 default-deny하며, kernel developer가 <code>vbpf_safe</code>로 승인한 경우만 허용한다. Pointer argument는 verifier의 type·memory-region 정보를 이용해 tenant-private/read-only 여부를 분류한다. 이 방식은 tenant code가 임의의 shared state mutation path를 우회하지 못한다는 invariant를 세우지만, annotation과 analyzer soundness를 trusted computing base에 포함한다.
# '''vBPF variables library'''
#: 복제 가능한 static state는 declarative API로 namespace별 instance를 lazy allocation한다. 여러 field와 이를 보호하는 lock은 한 storage group으로 배치해 일관되게 resolve한다. 반면 <code>sk_buff</code>처럼 복제할 수 없는 shared kernel object에는 [[BPF Type Format|BTF]] layout을 이용한 semantic state overlay를 적용한다. Helper 실행 전 clean base state를 복원하고 tenant patch를 apply하며, 실행 중 capture 또는 field-level update로 변경분을 기록한 뒤 namespace patch set에 finalize한다. Context switch 때 patch를 원자적으로 apply/restore하여 tenant별 view를 제공하며, full memory snapshot 대신 의미 있는 field 변경만 저장한다.
# '''Hot-path implementation'''
#: Sniffer registry와 Dispatcher table은 [[Read-copy-update|RCU]]-protected hash table로 구현하여 lookup을 lock-free로 수행한다. Critical-section allocation 실패와 latency variation을 줄이기 위해 object pool은 preallocate하고, tenant-private state는 첫 접근까지 lazy allocation한다. Hierarchy path와 BTF layout은 setup 시 precompute한다. 구현 규모는 Linux 6.12 kernel 약 12 KLoC와 LLVM 20 Clang plugin 약 1 KLoC이다.

== [[Conclusion]] ==
본 논문은 eBPF의 per-kernel context가 아니라 per-process context로 확장 시켜야 된다고 주장한다. 논문에서 제시하는 Motivation은<ref>Section 2.2 From Platform to Tenant eBPF</ref> 기존 논문이 다루지 않았다는 점에서 신선하였다. 그러나 본 논문은 필연적으로 커널의 여러 부분을 건드리기 때문에 여러 Corner case에 대한 의문이 생길 수 밖에 없었다.

우선, Motivation에서 궁금한 점이 있다. eBPF는 Per-process context로 실행시키는 프로그램이 있으면 새로운 BPF program type을 만들면 된다고 생각한다. XDP나 Storage같은 Device장치의 가속은 보통은 Per-process가 아니라 Per-kernel에서 작동해야만 효율적인 경우가 대부분인다. 따라서 eBPF의 현재 Execution model이 진짜 Per-process를 Support못해서가 아닌, 그렇게 하는 경우가 대부분인기 떄문에 지원하지 않는다고 보는 것이 좋아보인다. 그리고 eBPF는 기본적으로 CAP_BPF아니면 Root privileged를 요구하는 경우가 많다. 이 경우 단순히 그냥 namespace만 manual하게 체크할 수 있도록 하는 hook을 추가하면 될지 않을까 하는 생각이 든다.

Correctness측면에서의 의문은, 또한 Namespace로 묶는 vBPF Sniffer는 결국 Interrupt context -> task의 backward mapping인데, 이 부분이 모든 BPF Program type마다 다를텐데, 이를 어떻게 일이 태깅 하냐의 문제가 생길 것이라 생각한다. 또한 논문에서 주장하는 TCB에서 DoS가 빠져있는데, DoS는 Per-process context에서 매우 중요한 Threat model이다. 이를 뺴면 안된다고 생각한다. 또한 만약 Process에서 먼저 실행되지 않는 Context가 있으면 이는 어떻게 처리할 것인가? Kfunc는 이미 2022년 이후로 Default가 되었고, Helper function만 지원한다는 것도 큰 Limitation이다. Concurrency는 어떻게 할지도 의문이 남는다. RCU Lock을 잡고 Resource를 건들고 있다고 해보자. State overlay같은 경우에는 두개의 복사본이 있는 것인데, RCU Lock을 두 스레드가 공유하고 있다면? Merge할때 분명히 관련하여 문제가 생길것이라고 생각한다.

성능면에서는, 그리고 namespace lookup에 10ns-100ns가 걸린다고 하였는데, 이는 100Gb/s가 넘어가는 현재 시스템에 적용하기에는 큰 오버헤드이다.

최종적으로 여러 Per-tenent policy가 Global policy랑 Conflict될때 어떻게 Sound하게 처리할지에 대한 이야기가 없다. 만약 Global policy상 이 Process는 schedule out시켜야 하는데 per-tenent eBPF program이 schedule out시키지 않아야 한다고 주장하면 어떻게 할 것인가? Virtualization이 이 모든 Specific case를 보장할 수 있을 것인가? 어떤식으로 보장되는지 Restrict할수 있을 것인가?

[[분류: USENIX OSDI]]

Jwmalloc: A Verified Memory Allocator for Mobile Devices

2026-09-02T02:25:38Z

Noribot: Nori: update draft.wikitext

{{Paper
|title=jwmalloc: A Verified Memory Allocator for Mobile Devices
|author=Jiawei Wang, Ming Fu, Ruixian Wang, Chao Xu, Jonas Oberhauser, Haibo Chen
|conference=USENIX OSDI
|year=2026
}}

== 개요 ==
이 논문은 기존 범용 [[Memory allocator|동적 메모리 할당기]]의 설계가 모바일 workload의 잦은 memory reformatting, 큰 peak–trough memory swing, 공격적인 reclaim, 높은 thread oversubscription과 왜 맞지 않는지를 밝히고, uniform slab, closed sibling tree, lifetime-aware reclamation, non-blocking interface를 결합한 '''jwmalloc'''로 CPU·전력·메모리·tail latency를 함께 개선하였다.

== Motivation ==
모바일기기에서 Allocator의 성능은 사용가가 보는 응답에 직접 영향을 준다. 따라서 모바일 기기는 제한된 리소스인 CPU, Enenergy, Low memory overhead를 Low tail latency와 함께 제공해야 한다. 논문에서는 jemalloc이 안드로이드에서는 8.2%, 하모니 OS에서는 12.4%의 instruction cycles를 차지한다고 리포트 하였다.

저자들은 기존 allocator와 모바일 workload 사이에서 네 가지 mismatch가 있다고 주장한다.

# '''Frequent memory reformatting''' -> '''Size-class reuse overhead''': 시간에 따라 지배적인 object size가 빠르게 바뀌지만 전체 in-use memory는 비교적 일정하다. size class마다 slab 크기가 다른 jemalloc은 freed memory를 다른 class에 재사용할 때 backend subdivision/coalescing을 반복해야 한다.
# '''Large peak–trough swing''' -> '''Metadata overhead''': foreground/background 전환과 bursty interaction 때문에 peak memory가 steady-state보다 크게 높다. 관찰한 graphics service에서는 peak가 steady-state의 5배를 넘었다. peak에 맞춰 커진 grow-only metadata는 수요가 내려간 뒤에도 남는다.
# '''Aggressive reclamation''' -> '''System call overhead''': server에서는 5–15초인 reclaim delay가 모바일에서는 흔히 1초 이하이다. 너무 빨리 page를 OS에 돌려주면 곧 free될 이웃 range와의 coalescing 기회를 잃고 system call을 반복하지만, 무조건 늦추면 footprint가 증가한다.
# '''High oversubscription''' -> '''High parallelism requirement''': 한 app-market scenario에서 8 core 위에 123 process와 742 thread가 실행되었다. lock을 잡은 thread가 deschedule되면 UI thread까지 기다리는 [[Priority inversion]]이 생길 수 있으며, producer–consumer workload에서 흔한 cross-thread free가 이 위험을 키운다.

== Background ==
Memory Allocator는 대체로 small object를 slab에서 처리하는 frontend와, OS page 단위 virtual-memory range를 관리하며 slab 및 large object에 memory를 공급하는 backend로 나뉜다. jwmalloc은 4KB OS page 환경에서 2KB 미만을 thread-local frontend, 2–16KB를 process-wide slab-style midend, 16KB–4MB를 backend, 4MB 초과를 direct system call로 처리한다.

그러나, 모바일 allocation demand는 서버랑 다르게, 단순히 “작은 object가 많다”가 아니라, '''size별 demand가 빠르게 이동하고, page lifetime이 양극화되며, 적은 core에 많은 thread가 몰린다'''는 것이다.

따라서, 기존의 방식은 모바일 시스템에서 다음과 같은 4개의 문제가 있었다.

# CPU cost와 Slab관리로 인한 메모리 오버헤드의 Trade-off측면: Size-class를 관리하기 위한 CPU overhead가 크기 때문에 Frequent memory reformatting상황에서 CPU utilization이 커진다.
# OS에서 받아오는 Larger-size 단계의 존재: OS에서 메모리를 받아와서 각 size class별로 나누어 쓰기 때문에 Peak Memory상황에서 이 Two-level policy를 위한 Metadata의 Memory Overhead와 Division을 위한 CPU Utilization이 커진다.
# Memory reclamation policy가 Lifetime을 고려하지 않음: 이를 통해서 모바일 시스템에서는 Reclamation이 너무 빨라져서 추가적인 오버헤드가 발생함
# Coarse-grained lock을 통한 메타데이트 관리: High oversubscription이 자주 발생하는 모바일 환경에서는 적합하지 않은 디자인

== Main Idea ==
# Uniform and Pooled Slab Frontend: 모든 Slab이 같은 사이즈를 가지게 하였다.
# Size-class-exact range backend:
# Lifetime-based reclamation
# Non-blocking interfaces
# Verification under WMMs

== Design ==
; '''Uniform and Pooled Slab Frontend'''
: '''Local problem''': heterogeneous slab은 per-slab tail waste를 줄이지만 size demand가 바뀔 때 다른 크기의 slab로 즉시 바꿀 수 없어 backend churn을 만든다. 반대로 큰 uniform slab은 object 하나가 남아도 slab 전체가 pinned되는 footprint 문제가 있다.
: '''Mechanism''': half-page 이하 size class에 한 OS page(4KB)의 uniform slab을 사용하고, slab 크기를 바꾸는 대신 size class를 조절하여 tail fragmentation cap을 맞춘다. 완전히 빈 slab은 per-thread shared pool에 plain memory로 보관했다가 어느 frontend class로든 reinitialize한다. pool은 total slab count에 비례한 watermark와 LRU 정책으로 제한한다.
: '''Why it helps / tradeoff''': coalescing 없이 class 간 즉시 재사용할 수 있고 common-path instruction 수가 줄어든다. 다만 page에 가까운 size에서는 uniform slab의 이점이 약해지므로, 2–16KB object는 별도의 process-wide midend가 담당한다.
; '''Size-Class-Exact Range Backend'''
: '''Local problem''': fixed-granularity backend는 size class에 속하지 않아 request를 직접 만족하지 못하는 intermediate range를 만들고, eager coalescing은 곧 다시 split할 range까지 합쳐 CPU를 낭비한다.
: '''Mechanism''': larger range를 요청된 class와 나머지 valid class들로 '''eagerly subdivide'''하되, adjacent free range는 합친 결과 역시 valid size class일 때만 '''on demand coalesce'''한다. 각 root range는 resolution-3 '''closed sibling tree'''로 표현된다. 모든 node의 weight가 size-class 집합에 속하고, contiguous sibling들의 어떤 부분합도 다시 valid class가 되도록 tree를 구성한다.
: '''Why it helps / invariant''': backend가 유지하는 모든 range가 직접 사용할 수 있는 class 크기라는 invariant를 지켜 lookup과 반복 split/join을 줄인다. size-friendly multi-step split은 큰 contiguous remainder를 보존한다. 대신 binary buddy보다 tree construction과 concurrent state management가 복잡하다.
; '''Reclaimable Per-Granule Metadata'''
: '''Local problem''': fine-grained allocator의 per-range metadata는 빠른 direct indexing이 어렵고, 재사용되는 metadata record가 historical peak 이후 줄지 않을 수 있다.
: '''Mechanism''': closed sibling tree metadata를 4KB granule당 12B로 압축하고, 2차원 tree를 1차원 array에 mapping한다. internal-node 정보를 leaf에 encode하고 explicit parent/sibling pointer 대신 range size, depth, sibling index에서 address offset을 계산하는 '''metadata shifting'''을 사용한다.
: '''Why it helps / tradeoff''': direct lookup을 유지하면서 root range가 모두 free되면 metadata도 반환할 수 있다. 그러나 compact representation은 구현 추론을 어렵게 하며, 실제로 verification이 root node의 잘못된 sibling access bug를 발견했다.
; '''Lifetime-Based Reclamation'''
: '''Local problem''': fixed-delay reclaim은 곧 free될 이웃과 합칠 기회를 놓치거나, background 진입 후 allocator call이 줄어 reclaim 자체가 늦어질 수 있다. 모든 cached range의 timestamp를 scan하는 방식도 CPU가 든다.
: '''Mechanism''': mergeable sibling이 없는 free range는 countdown을 시작하고, sibling과 합쳐지면 이를 reset한다. mapped range의 size class마다 active/standby buffer를 두어 새로 free/coalesced된 range를 active에 넣는다. reclaim thread는 두 buffer를 swap하고 500ms를 기다린 뒤 standby의 physical page를 [[madvise]]로 버려 unmapped buffer로 옮긴다. cache가 watermark를 넘으면 freeing thread의 synchronous reclaim과 reclaim thread 조기 wakeup을 함께 사용한다.
: '''Why it helps / tradeoff''': long-lived allocation 옆의 free range부터 돌려주고 short-lived 이웃과 합쳐질 가능성이 있는 range에는 시간을 주며, per-range timestamp와 full scan을 피한다. 정책은 workload의 generational behavior와 interval/watermark 선택에 의존한다.
; '''Non-Blocking Interface and Ownership Transfer'''
: '''Local problem''': cross-thread free나 shared backend metadata의 coarse lock은 oversubscription에서 unbounded waiting과 UI stutter를 만들 수 있다.
: '''Mechanism''': 각 size class는 bounded concurrent bitmap을 fast path로 쓰고, bitmap이 비거나 가득 차면 locked unbounded list를 시도한다. lock 획득에 실패한 allocation은 더 큰 class의 range를 얻어 subdivide하고, free는 concurrent deferred list에 넣어 이후 operation이나 reclaim thread가 처리한다. Nest(search structure)와 Knit(contiguous-range tree)의 ownership은 staged acquire/release, rollback, atomic sibling-state update로 이전한다.
: '''Why it helps / tradeoff''': public allocation/free path가 lock holder를 기다리지 않고 진행하지만, rare contention을 extra subdivision, deferred work, 짧은 memory waste로 바꾼다. 따라서 non-blocking은 “비용이 없음”이 아니라 latency bound를 위해 일시적 resource overhead를 선택한 것이다.
; '''Bounded Verification under Weak Memory Models'''
: jwmalloc을 frontend/midend/backend의 library-style interface로 분리하고, cross-thread free, buffer drain, thread 생성·종료 등 edge case별 small concurrent client를 작성한다. [[VSync]] toolchain으로 각 client가 허용하는 weak-memory execution을 모두 탐색하여 assertion, memory safety, data-race absence, loop termination을 검사한다. state space를 감당하기 위해 bitmap entry 같은 constant는 logic을 유지한 채 축소했으며, 각 client는 10분 안에 완료되도록 구성했다.

== Result ==

=== Workload characterization ===

* Android와 HarmonyOS의 real-world workload에서 jemalloc instruction 비율은 평균 8.2%와 12.4%였다.
* graphics service는 burst 동안 약 3 million allocation operations/s에 도달했고, size별 activity는 시간에 따라 크게 변했다.
* streaming workload의 page 중 90%는 33.55ms 안에 free되지만 1% 이상은 3.22초 넘게 유지되었다. cross-thread free 비율은 대개 낮아도 짧은 구간에 80%, 일부 producer–consumer workload에서는 거의 100%에 접근했다.

이 결과는 jwmalloc의 세 정책 선택, 즉 cross-class slab reuse, lifetime-aware reclaim, non-blocking cross-thread path가 synthetic corner case만 겨냥한 것이 아니라 관찰된 모바일 behavior에 대응함을 보여준다.

=== Microbenchmarks ===

* 4-core x86 server에서 mobile oversubscription을 모사하고 rptest, xmalloc, mstress, mleak를 실행했다. jwmalloc은 jemalloc보다 평균 74% 높은 performance와 약 82% 적은 allocator-side instruction을 보였다.
* Frontend만 jwmalloc로 바꾸고 jemalloc backend를 유지한 ''jw+jemalloc''도 일부 test에서 크게 빨라져 uniform/pooling frontend의 독립적 효과를 보였다. Full jwmalloc의 rptest-8B-1MB-N에서는 jemalloc, mimalloc, tcmalloc의 allocator instruction이 각각 jwmalloc의 32.8×, 2.55×, 1.81×였다.
* Sleep-augmented mstress-10N에서 jwmalloc의 peak/steady footprint는 906MB/29MB였고 jemalloc은 968MB/376MB였다. ''jw+jemalloc''의 footprint가 jemalloc과 비슷했다는 점은 steady-state 감소의 주된 원인이 backend 및 reclamation임을 지지한다.
* Heavy oversubscription인 mstress-10N에서 jwmalloc의 P99.99 operation latency는 1.5µs로, 가장 좋은 경쟁 allocator의 5.9µs보다 낮았다.
* 단, allocation 분포가 4KB 근처에 몰린 8B–4KB/10N 설정에서는 uniform slab의 의도적 tradeoff와 fallback work 때문에 일부 performance/load regression이 있었다.

=== Real-world mobile scenarios ===

Huawei Mate 70 Pro(8-core ARMv8-A, 12GB RAM, HarmonyOS 5.1)에서 system jemalloc을 jwmalloc 또는 mimalloc로 교체하고, 32개 smartphone interaction으로 구성된 약 2시간 suite를 각 5회 실행했다.

* jwmalloc을 1로 normalize했을 때 jemalloc과 mimalloc의 전체 device instruction은 평균 1.10×와 1.13×였다. user-space allocation-related instruction은 각각 3.84×와 4.79×, kernel-space allocation-related instruction은 1.14×와 1.12×였다.
* Article reading과 video playback에서는 jemalloc 대비 instruction을 7–21%, CPU cluster power를 5–11%, LPDDR power를 2–3% 줄였다. 저자들은 CPU 절감을 lower voltage/frequency point 사용과 big core에서 small core로의 thread migration에 연결한다.
* 주요 system service의 [[Proportional set size|PSS]]는 대체로 jemalloc과 비슷하거나 낮았고, mimalloc은 흔히 더 많은 memory를 사용했다.

즉, microbenchmark의 allocator-side instruction 절감이 한 기기의 whole-system load와 CPU power 감소로 이어졌으며, 그 과정에서 jemalloc 대비 footprint가 악화되지 않았다는 것이 논문의 핵심 end-to-end evidence이다.

=== Verification and deployment ===

Bounded model checker는 개발 중 closed sibling tree의 root에 sibling이 있다고 잘못 가정하여 특정 data 값에서 out-of-bounds read와 crash를 일으키는 unknown bug를 약 10초 만에 찾았다. 이는 weak-memory-aware verification이 실제 구현 결함을 발견할 수 있음을 보이지만, 설정된 client와 bound 안의 보장이다.

저자들은 jwmalloc이 smartphone, tablet, smartwatch를 포함한 1,200만 대의 상용 기기에 배포되어 300억 user-hour 이상 안정적으로 동작했다고 보고한다. 이는 production maturity의 강한 관찰 증거이지만 formal correctness를 대체하지는 않는다.

== Contribution ==

# 모바일 allocation workload를 profile하여 frequent reformatting, peak–trough swing, aggressive reclamation, oversubscription이라는 네 가지 allocator mismatch를 정량화하였다.
# One-page uniform slab과 cross-size-class empty-slab pooling으로 frontend의 reformatting 비용을 줄였다.
# 임의 resolution-R size-class 집합을 exact하게 표현하는 closed sibling tree, eager subdivision/on-demand coalescing, reclaimable per-granule metadata를 제안하였다.
# Generational lifetime 관찰을 two-buffer tracker로 구현하여 timestamp/scan 없이 coalescing 기회와 timely reclaim을 조절하였다.
# Backend ownership protocol과 fallback/defer path를 통해 oversubscription에서도 allocation/free interface를 non-blocking으로 만들었다.
# Weak memory model 아래 bounded verification, microbenchmark ablation, real-device energy/footprint 평가, 대규모 production deployment를 함께 제시하였다.

== Conclusion ==
본 연구는 모바일 환경에서의 Unique한 특성을 Memory allocator에 반영시킬때 어떠한 점을 고려해야 하고, 어떠한 디자인을 적용시킬 수 있는지 제시하였다. 그 결과 Jemalloc대비 훨씬 더 좋은 성능을 보였다. Allocator만 최적화하여 10%단위의 성능 향상을 꾀하기에는 매우 힘들기에, 제시한 디자인의 최적화가 매우 효과적임을 시사한다.

=== Minor comments ===
# 섹션 2.2의 의도가 논문 전체에서 어떤 역활인지 모르겠다. Secure allocator가 모바일 환경에서 중요하지 않다는 정보가 갑자기 등장해서 굳이라는 생각이 들었다.

=== Major comments ===
# 기존의 Size-class방식이 더 잘 작동하는 Workload가 어떤 것인지 제시되어 있지 않다. Size-class로 나누는 이유는 Metadata management의 용이성 + Long-runnning시의 Fragmentation관리 때문인데, 본 시스템이 Size-class를 없앰으로서 어떤 Cons를 가지는지 분석한 결과가 있었으면 더 Complete했을 듯 하다.

[[분류: USENIX OSDI]]

PeeR: First-Class Scheduling for Latency-Critical eBPF Applications

2026-09-01T09:27:36Z

Noribot: Nori: update draft.wikitext

{{Paper
|title=PeeR: First-Class Scheduling for Latency-Critical eBPF Applications
|author=Jeremy Carin, Ben Holmes, Weiyang Wang, Ankit Bhardwaj, Manya Ghobadi
|conference=20th USENIX Symposium on Operating Systems Design and Implementation (OSDI 2026)
|year=2026
}}

== 개요 ==

이 논문은 latency-critical [[eBPF]] application이 non-preemptable [[softirq]]에서 실행될 때 발생하는 [[Head-of-line blocking]]과 CPU accounting 부재를, verifier가 보장하는 helper boundary에서의 [[Cooperative scheduling|cooperative preemption]]과 [[sched_ext]] 기반 2-level scheduling으로 해결하였다.

== Motivation ==
eBPF는 kernel 안에서 application logic을 실행하여 context switch와 data copy를 줄이지만, 그 실행 모델은 짧고 예측 가능한 packet filter를 가정한 채 run-to-completion으로 남아 있다<ref>While eBPF programs have grown dramatically in complexity, their execution model, which was originally designed for tiny packet filters, has remained largely unchanged.</ref>. 반면 최신 eBPF application은 복잡한 기능을 담당하는 Helper function을 eBPF program의 instruction대비 매우 많은 빈도로 실행시킨다. 이러한 환경에서는 eBPF Verifier의 eBPF 프로그램은 Instruction limit이 있기 떄문에 non-preemtible하다는 가정이 깨진다.

즉, 현재 복잡한 eBPF workload가 latency-critical I/O path의 softirq에서 실행되면 두 문제가 생긴다.

# '''Intra-application blocking''': long request가 먼저 도착하면 뒤의 short request가 preempt할 수 없어 [[Tail latency]]가 증가한다.
# '''Inter-application unfairness''': softirq 실행 시간은 eBPF application이 아니라 interrupt된 process에 귀속된다. 따라서 만약 Process A가 eBPF를 사용하게 된다면, eBPF 프로그램을 사용하지 않는 Process B도 Process A의 eBPF time을 paying해, scheduling에서 불리한 위치에 놓여지게 된다.

ksoftirq를 사용하더라도, ksoftirq는 invocation 내부를 preempt하거나 개별 eBPF task의 CPU 사용량과 우선순위를 scheduler에 드러내지 못한다. 따라서 application을 kernel fast path에 유지하면서도 invocation을 scheduling 가능한 resource로 만드는 새로운 Scheduling unit을 만들어야 한다.

== Importance ==
# PeeR는 eBPF의 안전성을 verifier 통과 여부에만 한정하지 않고, scheduler-visible accounting, preemption, policy enforcement까지 포함하는 resource-management 문제로 확장한다.
# [[eBPF verifier]]의 compile-time type/liveness 정보를 runtime continuation 복원에 재사용한다는 점이 핵심적인 architecture/runtime co-design이다.

== Challenges ==
eBPF를 schedulable하게 만드는 것은 다음의 요소 때문에 도전적인 과제이다.
# Enaling fine-grained preemption with low overhead
# Preserving eBPF correctness under preemption
# Providing expressive scheduling abstractions

== Main Idea ==
eBPF program이 kernel state와 상호작용하기 위해 helper function이나 kfunc를 자주 호출하고, verifier가 각 call boundary에서 register, stack, lock, reference의 상태를 이미 추적한다는 것이다. PeeR는 안전한 helper boundary에 JIT budget check를 삽입하여, budget을 넘긴 invocation이 다음 safe point에서 자동으로 yield하도록 만든다. 즉 application code에 명시적 <code>yield()</code>를 요구하지 않으면서 verifier-guided cooperative preemption을 구현한다.

실행 모델은 hybrid fast/slow path이다. budget 안에서 끝나는 task는 softirq에서 기존처럼 FCFS로 완료한다. budget을 넘긴 task만 상태를 continuation에 저장하고 per-CPU PeeR-kthread에서 복원한다. [[sched_ext]]가 outer macro-scheduler로서 eBPF workload와 userspace process 사이의 총 CPU 배분을 제어하고, kthread 내부의 micro-scheduler가 FIFO, [[Shortest remaining processing time|SRPT]], weighted round-robin 같은 policy로 개별 continuation의 순서를 정한다.

== Design ==

# '''Verifier-guided preemption point'''
#: '''문제''': timer interrupt로 임의 지점에서 끊으면 fast-path overhead가 크고, register·stack·kernel resource가 복원 가능한 상태인지 보장할 수 없다.
#: '''Mechanism''': PeeR는 JIT compiler가 안전한 helper/kfunc call 직전에 per-CPU epoch flag를 확인하는 prologue와 고유 <code>site_id</code>를 삽입하게 한다. budget이 만료되면 helper 실행 전에 yield handler로 분기하여 11개 register(88 B), 512 B eBPF stack, hook-specific context를 저장한다. flag가 꺼진 일반 경로의 check는 helper call당 1 cycle, scheduler가 flag를 갱신할 때는 2 cycle이다.
#: '''효과와 tradeoff''': 대부분의 짧은 task는 context switch 없이 기존 fast path를 유지하지만, preemption granularity는 안전한 helper 사이의 실제 실행 시간에 제한된다.
# '''Continuation 저장과 정확한 resume'''
#: '''문제''': softirq에서 저장한 stack/context pointer를 다른 kernel stack의 worker thread에서 그대로 복원하면 stale pointer가 되어 memory corruption을 일으킨다.
#: '''Mechanism''': verifier가 각 yield site에서 추적한 <code>PTR_TO_STACK</code>과 <code>PTR_TO_CTX</code> 정보를 patch descriptor로 보존하고, resume 시 새 stack/context base에 맞게 pointer를 rebase한다. JIT가 site별 [[Control-flow integrity|CFI]]-safe trampoline을 생성하므로, worker는 저장된 <code>site_id</code>에 대응하는 helper call 직전으로 돌아가 helper의 side effect를 정확히 한 번만 수행한다.
#: '''효과와 tradeoff''': source eBPF program 수정 없이 context를 옮길 수 있지만, verifier type이 모든 incoming path에서 일치하는 위치만 patch할 수 있다.
# '''Kernel resource와 [[Read-copy-update|RCU]] 안전성'''
#: '''문제''': softirq가 끝나면 implicit RCU read-side critical section도 끝나므로, live map-value/RCU pointer를 가진 채 yield하면 resume 전에 대상이 변경되거나 해제될 수 있다.
#: '''Mechanism''': verifier의 type 및 liveness 정보로 RCU-protected pointer, acquired reference, active lock이 없는 helper site에만 preemption check를 넣는다. lock-held region이나 unsafe resource가 live한 site는 yield point가 아니다.
#: '''효과와 tradeoff''': memory safety invariant를 보존하지만 cooperative gap이 길어질 수 있다. 또한 invocation이 서로 interleave되고 per-core completion order가 바뀔 수 있으므로 shared state의 원자성을 가정한 program은 spin lock 등으로 명시적으로 동기화해야 한다.
# '''Hook-specific continuation'''
#: '''문제''': verifier는 eBPF 내부 상태만 검증하므로, caller가 즉시 verdict를 요구하거나 lock을 쥔 hook은 결과를 나중에 적용할 수 없다.
#: '''Mechanism''': compatible hook은 input lifetime을 연장할 수 있고, caller가 deferred result를 허용하며, caller-side lock이 없어야 한다. 구현된 native XDP path는 packet을 reference-counted <code>xdp_frame</code>으로 전환하고 <code>XDP_YIELD</code>를 반환하여 [[NAPI]] poll을 계속한 뒤, worker에서 실제 DROP, ABORTED, PASS, TX, REDIRECT verdict를 적용한다.
#: '''효과와 tradeoff''': fast path가 deferred task를 기다리지 않지만 현재 구현은 native single-buffer XDP만 지원하며 multi-buffer와 tail-call program은 거부한다. 논문은 direct-action tc/TCX와 XRP storage hook이 model에 맞는다고 분석하지만 구현하지 않았다.
# '''Two-level scheduling과 accounting'''
#: '''문제''': sched_ext는 여러 continuation을 담은 하나의 PeeR-kthread만 보므로 개별 eBPF invocation의 순서나 owner별 CPU 시간을 직접 구분할 수 없다.
#: '''Mechanism''': outer sched_ext macro-scheduler는 <code>runtime_event</code> per-CPU ring buffer의 owner, CPU, duration, completion/yield 정보를 userspace daemon이 집계한 결과로 application별 CPU allocation을 조정한다. <code>set_budget()</code>은 기본 5 μs quantum을 설정한다. 각 per-CPU kthread 안의 micro-scheduler는 continuation priority와 <code>enqueue()</code>/<code>requeue()</code> callback으로 FIFO, SRPT, WRR을 구현한다.
#: '''효과와 tradeoff''': userspace와 eBPF를 하나의 CPU-sharing policy 아래 둘 수 있고 short/long request 사이의 policy도 선택할 수 있지만, task는 originating CPU에 고정되어 migration하지 않으며 SRPT에는 application-provided remaining-work estimate가 필요하다.

== Result ==
논문이 보고한 전체 범위에서 PeeR는 mixed workload의 latency-sensitive request p99를 current eBPF runtime보다 3–19.8배 줄였으며, 위 결과는 그 개선이 safe-point preemption과 scheduler-visible accounting이라는 두 mechanism에서 각각 발생함을 뒷받침한다.

== Contribution ==
# 7개의 modern eBPF application을 static analysis와 runtime profiling으로 조사하여, run-to-completion fast-path assumption이 execution-time variance, head-of-line blocking, CPU misaccounting 때문에 깨졌음을 정량화했다.
# verifier의 helper-boundary invariant와 type/liveness 정보를 이용해 low-overhead cooperative preemption, state capture, pointer rebasing, CFI-safe resume를 구현했다.
# softirq fast path와 per-CPU worker slow path를 결합하고, sched_ext macro-scheduler와 per-kthread micro-scheduler를 통해 intra-/inter-application policy를 함께 표현하는 2-level scheduling architecture를 제시했다.
# Redis, Memcached, synthetic TPC-C, colocation 및 microbenchmark에서 tail-latency 개선, CPU allocation enforcement, 247 ns preemption-resumption cost를 검증하고 quantum의 성능 tradeoff를 밝혔다.

== Conclusion ==
이 연구는 복잡해진 eBPF fast path를 단순한 안전한 kernel extension이 아니라 CPU accounting과 scheduling이 필요한 first-class workload로 바라보게 만든다. PeeR는 verifier가 보장하는 helper boundary를 safe point로 재사용하고, 짧은 task는 softirq에 남기되 긴 task만 continuation으로 전환하여 low overhead preemption과 policy control을 양립시킬 수 있음을 보였다. 다만, Correctness측면에서 PeeR을 Support하기 위한 방법이 XDP가 아니라 좀더 복잡한 영역으로 확장 가능할지는 의문이다. 만약 할당된 Struct를 eBPF program내부에서 수정하고, Helper function에 의해서 Preemption된다면, 기본적으로 atomic을 가정하는 Helper function입장에서는 [[TOCTOU]]가 발생할 수 있다. 또한 전체 eBPF의 Kfunc라던지 이러한 시스템에 통합하기 위해서는 많은 Kernel-patch가 필요한 점도 Implementation측면에서의 약점이 될 수 있을듯 하다. 그러나 Motivation이 매우 좋고, Design또한 Motivation에 잘 이어지며, 전체적인 시스템의 Nolvety도 좋아보인다. 그리고 Implementation details들이 잘 나와 있어서, Sound하다는 인상을 받았다.

[[분류: USENIX OSDI]]

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]]