CuFuzz: Hardening CUDA Programs through Transformation and Fuzzing

CuFuzz: Hardening CUDA Programs through Transformation and Fuzzing
Author	Saurabh Singh, Ruobing Han, Jaewon Lee, Seonjin Na, Yonghae Kim, Taesoo Kim, Hyesoon Kim
Conference	arXiv:2601.01048v1 [cs.CR]
Year	2026

개요

이 논문은 CUDA 프로그램이 GPU의 대규모 병렬 실행, 복잡한 메모리 계층, 부정확한 오류 보고 때문에 기존 CPU용 Fuzzing 도구의 혜택을 받기 어려운 문제를 다루며, GPU 코드를 CPU 코드로 변환하되 보안 버그와 관련된 메모리 접근을 보존하여 AFL++와 AddressSanitizer를 재사용하는 compiler-runtime co-design 도구 CuFuzz를 제안하였다.

Background

CUDA는 C/C++의 저수준 제어와 성능을 계승하지만 Buffer overflow, Out-of-bounds access, Use-after-free 같은 메모리 안전성 문제도 함께 계승한다. GPU에서는 global, local, shared memory가 서로 다른 수명과 가시성을 가지며, 수천 개의 thread가 SPMD 방식으로 실행된다. 또한 CPU의 segmentation fault와 같은 precise exception이 부족하여 잘못된 접근이 즉시 crash나 정확한 오류로 이어지지 않을 수 있다.

Motivation

Coverage-guided fuzzing은 CPU 생태계에서 AFL++, AddressSanitizer와 결합해 대량의 입력을 실행하고 crash를 정확한 피드백으로 사용할 수 있다. 반면 GPU는 다음 이유로 같은 방법을 직접 적용하기 어렵다.

대규모 thread/block hierarchy를 그대로 CPU에서 실행하면 한 입력의 실행 비용이 매우 커진다.
global, local, shared memory와 CUDA 전용 API를 CPU의 메모리·실행 모델에 대응시켜야 한다.
GPU의 오류 검출과 보고는 CPU보다 제한적이어서 fuzzing의 핵심 피드백인 crash를 놓칠 수 있다.
전용 GPU fuzzing infrastructure를 새로 구축하면 기존 CPU fuzzer, sanitizer, 서버 자원을 재사용하기 어렵다.

따라서 GPU의 수치 출력을 완전히 재현하는 범용 migration보다, 메모리 버그를 드러내는 접근 행동을 보존하면서 높은 exec/sec를 얻는 fuzzing 전용 변환이 필요하다.

Importance

이 논문의 중요한 전환은 GPU-native detector를 새로 설계하는 대신, CUDA kernel의 bug-relevant behavior를 CPU 실행으로 옮겨 성숙한 CPU fuzzing·sanitizer 생태계를 활용한다는 점이다. 특히 fuzzing에는 완전한 numerical equivalence가 아니라 메모리 버그 보존이 핵심이라는 관찰을 optimization boundary로 삼는다.

또한 저자들은 global/local/shared memory의 spatial·temporal 오류를 분류한 100개 test의 GMSBench를 제시한다. 이 분류와 도구별 coverage 비교는 CuFuzz 자체뿐 아니라 향후 GPU memory-safety detector를 비교하는 evaluation framework로도 재사용할 수 있다.

Main Idea

핵심 아이디어는 CUDA host/kernel code를 LLVM IR 수준에서 CPU-executable program으로 변환하고, 원래 GPU program의 메모리 공간과 접근을 CPU의 heap/stack 접근으로 대응시킨 뒤 AFL++와 CPU memory-error detector로 실행하는 것이다. 이때 kernel의 수치 결과까지 동일하게 유지하는 대신, memory access instruction과 그 index에 영향을 주는 계산을 보존하는 것을 fuzzing correctness의 기준으로 삼는다.

변환만으로는 GPU의 수많은 thread를 CPU loop로 실행하는 비용이 너무 크다. CuFuzz는 affine access kernel에서 경계 thread/block만 우선 실행하는 Partial Representative Execution (PREX)과, memory-access index에 영향을 주지 않는 계산과 barrier를 제거하는 Access-Index Preserving Pruning (AXIPrune)으로 throughput을 높인다.

Design

cuf-cc compilation pipeline
- Local problem: CUDA executable은 CPU에서 실행되는 host code와 GPU용 NVVM IR로 내려가는 kernel code를 함께 포함하므로 일반 CPU fuzzer가 kernel을 직접 instrument할 수 없다.
- Mechanism: nvcc/clang의 drop-in replacement인 cuf-cc가 host-side LLVM IR와 kernel-side NVVM IR를 분리한 뒤 둘 다 x86_64 LLVM IR로 변환한다. Kernel function에는 AddressSanitizer annotation을 추가하고 AFL++의 branch-coverage instrumentation을 적용한다.
- Effect and tradeoff: 기존 CPU fuzzing toolchain을 재사용할 수 있지만, 검출 가능한 버그의 범위는 선택한 CPU detector와 변환이 보존하는 semantics에 의존한다.
PREX-Aware Collapsing Transform (PACT)
- Local problem: GPU의 gridSize × blockSize thread를 CPU thread로 일대일 대응시키면 실행 비용이 지나치게 크다.
- Mechanism: GPU block 하나를 CPU thread 하나에 매핑하고, block 내부 GPU thread는 loop iteration으로 직렬화한다. __syncthreads() 전후는 별도 loop로 분할하며, barrier를 가로질러 살아 있는 local variable은 array로 확장한다. GPU global memory는 CPU heap으로, shared/local memory는 CPU stack으로 대응시킨다.
- Invariant and tradeoff: 각 phase의 memory ordering과 memory-space별 overflow의 가시성을 보존하지만, GPU의 병렬 interleaving보다 더 강한 순서를 도입하고 barrier마다 loop와 array overhead가 생긴다.
Partial Representative Execution (PREX)
- Local problem: 변환 후에도 모든 block과 thread를 매 입력마다 실행하면 fuzzing throughput이 낮다.
- Mechanism: compiler의 Affine Access Analysis가 memory index를 thread/block index의 affine function으로 표현할 수 있는지 판정한다. 조건문 밖의 affine access에 대해서는 첫/마지막 block의 첫/마지막 thread가 경계 OOB를 대표한다는 성질을 이용한다. 조건문이 있는 경우 runtime coverage monitor가 바깥쪽 block부터 점진적으로 추가하며 memory-access coverage가 완료되면 중단한다.
- Invariant and tradeoff: affine-access kernel에서는 representative execution으로 작업량을 줄이고, non-affine kernel에서는 모든 block/thread 실행으로 fallback한다. 효율은 affine 판정과 coverage가 bug-relevant behavior를 충분히 대표한다는 가정에 달려 있다.
Access-Index Preserving Pruning (AXIPrune)
- Local problem: expensive math와 barrier가 CPU 변환 뒤 많은 계산, loop, auxiliary memory access를 만들지만 상당수는 memory bug 발생 위치와 무관하다.
- Mechanism: data-dependence analysis로 branch condition이나 memory-access index에 영향을 주는 값을 추적한다. 관련 없는 math operation을 입력값 전달로 대체하고, shared-memory value가 branch/index에 영향을 주지 않으면 barrier를 제거한다.
- Invariant and tradeoff: 원래 GPU program의 memory access와 index 관련 계산은 남기지만 numerical output equivalence는 의도적으로 포기한다. 따라서 일반 CUDA-to-CPU translator가 아니라 memory-bug fuzzing 전용 optimization이다.
CuFuzz runtime
- CUDA allocation/copy API의 CPU 구현과 thread pool을 제공하고, kernel launch를 번역된 function call로 바꾼다. Runtime library도 AddressSanitizer로 instrument한다. PREX mutator는 각 fuzzing iteration에서 실행할 block을 선택하고 coverage를 감시한다.

Result

Error-detection coverage

저자들은 100개 test로 구성된 GMSBench에서 native CUDA와 여러 GPU/CPU detector의 검출 범위를 비교한다.

실행/검출 조합	검출 수	Coverage	해석
Native CUDA	8/100	8%	별도 보호가 없을 때 allocator가 보고하는 일부 temporal error에 한정된다.
NVIDIA Compute Sanitizer	36/100	36%	tripwire 방식이므로 buffer 경계를 넘어도 다른 유효 영역에 머무는 arbitrary read/write 등을 놓친다.
CuFuzz + AddressSanitizer	48/100	48%	host와 kernel을 함께 검사할 수 있고 intra-allocation 일부를 찾지만, tripwire와 allocator 차이의 한계가 있다.
CuFuzz + No-Fat	91/100	91%	exact base/bounds에 가까운 검사를 통해 가장 높은 coverage를 보이지만 dynamic shared-memory와 일부 allocator-mismatch case는 남는다.

이는 GPU-to-CPU translation 자체보다 뒤에 결합하는 detector가 최종 bug coverage를 크게 좌우함을 보여 준다. 논문의 주 fuzzing campaign은 coverage가 더 높은 No-Fat이 아니라 AddressSanitizer를 사용했다.

Fuzzing campaign

HeteroMark, Rodinia, CUDA SDK에서 고른 14개 benchmark를 각각 12시간 동안 AFL++로 fuzzing했다. 각 benchmark에는 별도 fine-tuning 없이 seed file 하나를 주었고, AFLTriage로 결과를 deduplicate하였다. CuFuzz + AddressSanitizer는 총 122개 finding을 보고했으며 구성은 kernel crash 15개, host crash 40개, hang 67개이다.

Kernel에서는 matmul, bfs, cfd, lud의 global-memory OOB read/write를, host에서는 aes, ga, pr, bfs의 heap-buffer overflow를 찾았다. cfd의 null-pointer dereference, ga와 ep에서 memcpy()에 null pointer가 전달되는 경우, 여러 benchmark의 invalid allocation·integer overflow·argument-parser exception도 보고했다. 또한 malformed input의 0 값이 task-count 계산에서 division-by-zero를 일으키는 CuFuzz runtime 자체의 bug도 발견했다. 이 결과는 변환된 실행이 host와 kernel 양쪽의 crash/hang triage에 실용적인 피드백을 줄 수 있음을 보이지만, 122개 모두가 독립적으로 검증된 보안 취약점이라는 뜻은 아니다.

Throughput

100회 실행의 평균 exec/sec로 측정했을 때 PREX는 naïve translated baseline 대비 14개 중 12개 benchmark에서 성능을 높였고 평균 약 32×, affine access가 크고 thread 수가 많은 workload에서 최대 183× 향상을 보였다. pr과 bfs는 indirect access 때문에 non-affine으로 판정되어 partial execution의 이득이 없었다.

PREX 적용 코드 위에서 AXIPrune은 평균 33%의 추가 throughput 향상을 냈다. 예를 들어 ep에서는 memory index와 무관한 exponential function 제거로 27% 향상되며, hist에서는 index와 무관한 barrier 및 변환이 만든 보조 접근을 줄인다. 두 optimization을 합친 최대 speedup은 naïve translation 대비 224.31×이다. 즉 PREX는 실행할 GPU thread/block 수를 줄이고 AXIPrune은 각 representative execution의 불필요한 instruction을 줄여 서로 다른 병목을 완화한다.

Contribution

CUDA program을 CPU-executable LLVM IR로 변환하고 AFL++·CPU memory detector와 연결하는 end-to-end compiler-runtime co-design fuzzing framework를 제안하였다.
Affine memory-access pattern을 이용해 representative block/thread를 선택하는 PREX와 access-index dependency를 보존하며 불필요한 연산을 제거하는 AXIPrune을 제안하였다.
GPU global/local/shared memory의 spatial·temporal bug를 체계화한 100-test GMSBench와 detector coverage 비교를 제공하였다.
14개 CUDA benchmark의 12시간 campaign에서 122개 deduplicated crash/hang finding을 얻고, naïve translation 대비 평균 약 32× 및 최대 224.31×의 throughput 향상을 보였다.

Related Work Position

CuFuzz는 AFL++, AddressSanitizer, libFuzzer 같은 CPU fuzzing infrastructure를 GPU program에 연결한다는 점에서 CPU-only fuzzer와 다르다. 기존 OpenCL kernel test generation은 mutation과 selective constraint solving을 결합했지만, CuFuzz는 CUDA 전체 program의 host/kernel translation과 sanitizer integration을 목표로 한다.

GPU memory-safety 측면에서 NVIDIA Compute Sanitizer, GPUShield, CuCatch, GMOD는 GPU에서 직접 오류를 검출하거나 bounds protection을 제공한다. CuFuzz는 방어 메커니즘을 GPU에 배치하기보다 실행을 CPU로 옮겨 여러 CPU detector를 교체해 사용할 수 있게 한다. CuPBoP과 같은 GPU-to-CPU migration tool은 numerical correctness를 우선하지만, CuFuzz는 debugging symbol, fuzzing instrumentation, bug-oriented partial execution/pruning에 초점을 둔다.

Criticisms

변환이 보존하는 버그의 범위: PACT는 block 내부 thread를 loop로 직렬화하고 CPU heap/stack 및 allocator semantics를 사용한다. 따라서 data race, warp scheduling, weak memory ordering, asynchronous API behavior, GPU hardware/driver 고유 오류처럼 concurrency나 device semantics에 의존하는 취약점은 제거되거나 새롭게 생길 수 있다. 논문은 주로 memory-access safety를 다루며 이러한 bug class에 대한 preservation을 실험하지 않는다.
PREX·AXIPrune의 가정: PREX의 강한 대표성 정리는 affine index와 조건문 밖의 access를 전제로 하고, 조건문은 runtime coverage로 완화한다. 그러나 memory-access instruction coverage가 모든 동적 index/state interaction을 대표하는지, AXIPrune의 dependency analysis가 pointer aliasing이나 복잡한 shared-memory dataflow에서도 bug-triggering behavior를 항상 보존하는지는 더 강한 검증이 필요하다.
제한된 campaign: 평가는 세 benchmark suite의 14개 program, benchmark당 12시간, 한 hardware/software configuration에 한정된다. Production CUDA application, 반복 trial, variance, time-to-first-bug, coverage growth curve가 없어 대규모 실제 서비스에 대한 일반화와 fuzzing 효율의 통계적 안정성을 판단하기 어렵다.
Finding과 vulnerability의 구분: 122개 finding에는 67개 hang과 malformed-input에 의한 allocation failure, exception, runtime bug가 포함된다. Patch, maintainer confirmation, exploitability analysis, CVE assignment가 제시되지 않으므로 이를 122개의 독립적·확정된 security vulnerability로 해석해서는 안 된다.
Detector 및 baseline 비교: GMSBench에서 CuFuzz + AddressSanitizer coverage는 48%이고 CuFuzz + No-Fat은 91%이지만 실제 fuzzing campaign은 AddressSanitizer만 사용한다. 또한 GMOD, GPUShield, CuCatch, No-Fat의 일부 결과는 구현을 직접 실행한 것이 아니라 해당 논문의 설명을 바탕으로 추정했다. Performance baseline도 naïve CPU translation이므로 native GPU tool이나 다른 end-to-end GPU testing system과의 bug-finding cost 비교가 부족하다.
재현성: 논문은 GMSBench를 공개할 예정이라고 표현하지만, 제공된 preprint만으로는 artifact availability, source revision, 전체 실행 script를 확인할 수 없다.

Conclusion

이 연구는 CUDA fuzzing을 GPU-native 오류 보고의 문제로만 보지 않고, bug-relevant memory behavior를 CPU로 옮겨 성숙한 fuzzing·sanitizer infrastructure를 재사용하는 문제로 재구성한다. CuFuzz는 PACT, PREX, AXIPrune을 통해 이 변환의 실행 비용을 크게 줄이고 benchmark campaign에서 host/kernel crash와 hang을 실제로 수집할 수 있음을 보였다. 향후 GPU fuzzing 연구에서 기억할 핵심은 numerical equivalence보다 어떤 bug semantics를 보존할 것인지 명시하고, 그 보존 범위와 detector coverage를 함께 평가해야 한다는 점이다.