---
title: 'cuFuzz: Multi-Use Fuzzing Framework'
url: https://www.emergentmind.com/topics/cufuzz-43bf24bb-0e51-4990-a245-3f3954ad9c4d
type: topic
---

# cuFuzz: Multi-Use Fuzzing Framework

cuFuzz is a name used for several distinct fuzzing artifacts in recent literature. Most prominently, it denotes two 2026 CUDA-focused systems: a CUDA-oriented, whole-program, coverage-guided fuzzer that couples AFL++ with NVBit-based device instrumentation and separate sanitizer processes, and a compiler-runtime co-design framework that transforms CUDA programs into CPU programs so that AFL++ and AddressSanitizer can be applied to bug-preserving CPU executions. In a separate 2023 usage, cuFuzz denotes a fork-awareness evaluation idea for coverage-guided fuzzers rather than a standalone implementation [2603.12485][2601.01048][2301.05060].

## 1. Nomenclature and scope

The term is not monosemous. Two papers use **cuFuzz/CuFuzz** as the name of a CUDA fuzzing system, while another source explicitly states that, in its context, **cuFuzz is not a specific fuzzer implementation** but an evaluation framework for fork-awareness. The supplied summary of "Fuzzing with Quantitative and Adaptive Hot-Bytes Identification" further notes that the paper’s actual system name is **Finch**, even though the explanation refers to it as **cuFuzz/Finch** because of the query wording [2603.12485][2601.01048][2301.05060][2307.02289].

| Usage | Technical description | Paper |
|---|---|---|
| cuFuzz | CUDA-oriented, whole-program, coverage-guided fuzzer | [2603.12485] |
| CuFuzz | Compiler-runtime co-design framework transforming CUDA to CPU for fuzzing | [2601.01048] |
| cuFuzz | Fork-awareness evaluation idea for coverage-guided fuzzers | [2301.05060] |
| Finch | Quantitative and adaptive hot-bytes identification; actual system name is Finch | [2307.02289] |

This naming collision is consequential for technical reading. Results, architectures, and threat models attached to one usage cannot be transferred to the others without qualification.

## 2. Whole-program CUDA fuzzing

"Hunting CUDA Bugs at Scale with cuFuzz" presents cuFuzz as a **CUDA-oriented fuzzer** whose central design decision is **whole program fuzzing** rather than kernel-level fuzzing. The paper identifies three obstacles in prior GPU fuzzing efforts: **kernel-level fuzzing leading to false positives**, **lack of device-side coverage-guided feedback**, and **incompatibility between coverage and sanitization tools**. cuFuzz addresses these by using the program’s real entry point as the harness: for open-source programs it fuzzes the application’s `main`, and for closed-source CUDA libraries it uses supplied sample applications as harnesses. This avoids impossible host-device states such as null pointers or out-of-range sizes that would never survive host-side validation, and therefore avoids the false positives associated with independently fuzzing device-side kernels [2603.12485].

The emphasis on whole-program execution is significant because CUDA execution is heterogeneous by construction. Host code determines kernel selection, launch parameters, and preconditions. The paper’s examples show that device failures may be downstream consequences of host-side failures, including unchecked CUDA API errors. Under that model, bug discovery is not reducible to device-kernel mutation alone.

## 3. Unified host-device feedback and sanitization separation

cuFuzz combines **host-side** and **device-side** coverage into a single feedback loop. Host-side coverage is collected with ordinary **AFL++ compile-time instrumentation**. Device-side coverage is collected at runtime with **NVBit**, which patches GPU basic blocks at load time and lets cuFuzz record edge transitions from device execution. The host-side edge tracker is AFL-style:
$$
idx = \_\_afl\_prev\_loc \oplus cur\_loc,
$$
with the bitmap counter incremented at `__afl_area_ptr[idx]` and `__afl_prev_loc = cur_loc >> 1`. cuFuzz extends this idea to the GPU and maps device coverage into the same bitmap while keeping namespaces disjoint. In the prototype, the bitmap is **64 KB**; the first **32 KB** are reserved for host edges and the second **32 KB** for device edges. The device-side index is
$$
idx = (MAP\_SIZE/2) + ((prev\_loc[tid] \oplus cur\_loc) \bmod (MAP\_SIZE/2)).
$$
Because many GPU threads can execute the same branch simultaneously, cuFuzz uses **warp-aware atomic updates** and **power-of-two buckets**; for each warp, only the first active lane performs the atomic increment. An input is retained if it discovers any new host edge or any new device edge, so mutation is steered by a genuinely unified host-device feedback channel [2603.12485].

A second architectural contribution is the decoupling of **sanitization** from **coverage collection**. The paper states that CUDA sanitizers and GPU binary instrumentation tools conflict through shared runtime dependencies. cuFuzz therefore runs the coverage-guided harness in one process and launches sanitizer-enabled wrappers in separate processes inside the fuzzing loop. The sanitizers include **AddressSanitizer** for host-side memory bugs and NVIDIA Compute Sanitizer’s **`memcheck`**, **`racecheck`**, and **`initcheck`** for device-side memory errors, races, and uninitialized reads. The implementation uses **selective sanitization**, inspired by **SAND**, so only selected inputs such as those that reach new coverage are sent to the sanitizer processes. The implementation is built on **AFL++ 4.31c**, **NVBit 1.7.5**, **NVIDIA Compute Sanitizer 2025.2.0.0**, and **CUDA 12.9**.

Persistent-mode support is realized by modifying the harness to run multiple test cases per process using AFL++ loop mode. Because NVBit must know iteration boundaries, the harness launches a tiny **notification kernel** with a magic ID at the beginning and end of each loop iteration; NVBit intercepts those calls, resets device coverage at iteration start, and merges host/device coverage at iteration end. The harness also writes each testcase to a file so the separate sanitizer processes can consume the same input.

## 4. Empirical performance and bug discovery in CUDA workloads

The evaluation of the whole-program cuFuzz spans **14 CUDA workloads** drawn from **HeCBench** and **NVIDIA CUDA library samples**, including machine learning, math, image processing, compression, and random number generation. Across **three 24-hour runs per configuration**, cuFuzz found **43 previously unknown bugs**, including **19 in commercial libraries**. The paper reports that all bugs were acknowledged by maintainers and that **40 of the 43** had already been fixed in upstream releases by the time of writing. The bug distribution comprises **8 heap buffer overflows**, **1 stack buffer overflow**, **2 floating-point exceptions**, **2 segmentation faults**, **13 out-of-bounds device accesses**, **5 shared-memory races**, and **12 uninitialized device reads** [2603.12485].

The results are especially strong on closed-source targets. For open-source workloads, plain AFL++ often performs similarly or even better in raw coverage because it runs faster and host coverage is often sufficient. For closed-source libraries, however, host coverage alone is insufficient, and device-side instrumentation becomes decisive. The paper reports device-side edge improvements over a no-device-coverage variant ranging from **9% for nvTIFF** to **289% for blas-gemm**. It also states that cuFuzz reaches the same coverage as AFL++ with fewer executions because the additional GPU feedback makes mutations more informative.

The overhead analysis is explicit. Device-side coverage collection with NVBit is the largest throughput cost, reducing throughput by **67%** for device-triggering inputs. Host-side coverage and AddressSanitizer have **negligible throughput impact**. Device-side sanitizers also impose substantial costs: **memcheck** reduces throughput by **39%**, **racecheck** by **66%**, and **initcheck** by **32%**. These measurements motivate both decoupled sanitization and selective sanitization.

Persistent mode improves performance for **50% of the benchmarks**, finds **35 of the 43 bugs**, and discovers **16 bugs faster than any other configuration**. In **11 of 14 programs**, persistent mode is equal or better than single-process-per-input execution, although it does not help universally. The paper attributes the negative cases to reinitialization costs for large GPU allocations and to very short-lived kernels where loop bookkeeping offsets the gain.

A head-to-head comparison with kernel-level fuzzing on **nine open-source benchmarks** is one of the sharpest results in the paper. Kernel-level fuzzing found only **6 of 14 device-side bugs**, missed **8** because of violated host-enforced invariants, and produced **16 false positives**. cuFuzz found **all 14 device-side bugs** plus **10 additional host-side bugs** with **zero false positives**. The paper’s limitation analysis nevertheless notes that throughput remains much lower than in CPU fuzzing and that the current coverage model does not distinguish thread interleavings, so race bugs dependent on scheduling may still be missed.

## 5. Translation-based CuFuzz

"CuFuzz: Hardening CUDA Programs through Transformation and Fuzzing" defines CuFuzz as a **compiler-runtime co-design framework** that brings **coverage-guided fuzzing to CUDA programs** by **transforming GPU code into CPU code** and then reusing mature CPU fuzzing infrastructure such as **AFL++** and **AddressSanitizer**. Its core thesis is not native device instrumentation but **translation**: make CUDA programs look like CPU programs while preserving the memory behavior relevant to bug finding. The motivation is that **CUDA inherits C/C++ memory unsafety**, and GPU programs may contain buffer overflows, out-of-bounds reads and writes, use-after-free or use-after-scope, invalid free or double free, and uninitialized memory issues; prior work, as summarized by the paper, shows that such bugs can become exploit primitives including data corruption, control-flow hijacking, denial of service, and information leakage [2601.01048].

The pipeline has two stages. In the **compilation stage**, `cuf-cc` acts as a drop-in replacement for `nvcc` or `clang` and compiles CUDA into LLVM IR for both kernel/device code as **NVVM IR** and host code as **x86 LLVM IR**. A kernel translator then turns GPU IR into CPU-targeted LLVM IR. In the **fuzzing stage**, the translated IR is optimized, **ASan instrumentation** is inserted, **AFL++** instruments branch coverage, and the translated binary is fuzzed like an ordinary CPU program. CuFuzz also provides runtime libraries that emulate CUDA functions such as `cudaMalloc()`, `cudaMemcpy()`, and kernel launch handling, and these runtime libraries are themselves ASan-instrumented.

The transformations are explicit. On the host side, the translator replaces CUDA runtime calls with CPU equivalents, converts kernel launch syntax into direct function calls, removes calls such as `cudaDeviceSynchronize()` that are unnecessary in CPU execution, and maps CUDA constant-memory accesses to regular host memory operations. On the kernel side, it converts GPU thread hierarchy into CPU loops or threads, remaps **global memory to heap**, **shared memory and local memory to stack**, and aims to preserve memory bugs by keeping the relevant access patterns intact.

A key execution transform is **PACT**, the PREX-aware collapsing transform. Under PACT, each GPU **block** is mapped to **one CPU thread**, and threads inside the block are mapped to **loop iterations**. If the kernel contains `__syncthreads()`, PACT splits execution into separate loops before and after the barrier, and local variables used across the barrier are promoted into arrays to preserve correctness. The paper explicitly notes that this is semantically conservative: the CPU version may impose stricter ordering than the GPU version, but it preserves the bug-relevant memory behavior.

## 6. PREX and AXIPrune

The translation-based CuFuzz introduces two throughput-oriented co-optimizations: **Partial Representative Execution (PREX)** and **Access-Index Preserving Pruning (AXIPrune)**. PREX exploits the fact that many CUDA kernels are **SPMD** and that, for **affine-access kernels**, memory access indices can be written as
$$
\overrightarrow{indices} = \mathbf{T}
\begin{bmatrix}
\text{threadId} \\
\text{blockId}
\end{bmatrix}
+
\overrightarrow{b},
$$
where $\mathbf{T}$ and $\overrightarrow{b}$ are invariant across threads. The paper states that if a memory bug exists, then one of the boundary threads is sufficient to expose it:
$$
bug\_threads \cap \{t_0^0, t_0^{B-1}, t_{T-1}^0, t_{T-1}^{B-1}\} \neq \varnothing.
$$
PREX therefore executes representative boundary blocks and threads, beginning with the first and last blocks and stopping early if coverage reaches **100%** or a bug is found. The reported effect is an **average 32× speedup**, **up to 183× speedup**, with throughput improvements in **12 out of 14 benchmarks** [2601.01048].

AXIPrune removes instructions that do **not affect memory access indices**. Its two main components are **barrier elimination** and **math elimination**. Barriers such as `__syncthreads()` are removed when the shared-memory values involved are not used in branch conditions or memory access indices. Expensive math instructions are removed when their outputs do not affect branches or memory indices. The paper is explicit that AXIPrune is **not intended to preserve full numerical semantics**; it preserves memory-access behavior needed for fuzzing. On top of PREX-optimized code, AXIPrune yields an **additional 33% throughput gain on average**, and together the two optimizations can reach **up to 224.31× speedup**.

The evaluation covers **14 CUDA benchmarks** from **HeteroMark**, **Rodinia**, and **CUDA SDK samples**, and introduces **GMSBench**, an open-source CUDA memory-safety benchmark suite with **100 test cases**. CuFuzz uncovered **122 unique issues** in its fuzzing campaigns, including **15 kernel crashes**, **40 host crashes**, and **67 hangs**; after AFLTriage triage, the paper reports **55 crashes total**, **67 hangs**, and **122 total findings**. Reported bug classes include kernel global memory out-of-bounds reads and writes, host heap-buffer overflows, null pointer dereferences, null passed to `memcpy()`, unhandled exceptions in argument parsing, division-by-zero bugs in runtime task calculation, malformed-input-triggered crashes, and bugs in the CuFuzz runtime itself. The paper’s stated limitations are that PREX applies only when access patterns are affine, AXIPrune trades semantic fidelity for fuzzing speed, and the framework is **bug-preserving, not output-equivalent**.

## 7. Fork-awareness usage and terminological ambiguity

In "Evaluating the Fork-Awareness of Coverage-Guided Fuzzers," **cuFuzz is not a specific fuzzer implementation**. It denotes a **fork-awareness evaluation idea** for coverage-guided fuzzers, specifically a way to assess whether fuzzers can properly handle programs that use `fork()`. The paper defines fork-awareness as follows: **"A coverage-guided fuzzer is fork-aware if it can detect bugs and hangs and measure coverage in the same way for both the child and the parent's branch."** The property is decomposed into **C1: Child bugs detection**, **C2: Child hangs detection**, and **C3: Child code coverage**. The evaluation covers **14 coverage-guided fuzzers**, including **11 from FuzzBench** and **3 from ProFuzzBench**, using three small C challenge programs for child-process crash detection, child-process hang detection, and child-branch coverage measurement. The headline findings are **0/14** satisfying all fork-awareness criteria, **11/14** failing to detect child-process bugs, **0/14** detecting child-process hangs, and **14/14** measuring child coverage in some form [2301.05060].

This usage differs categorically from the CUDA systems. It is a property-based evaluation framework centered on `fork()` behavior, not a GPU fuzzer, not a compiler transformation tool, and not a CUDA runtime instrumentation system.

A further source of ambiguity appears in the supplied summary of "Fuzzing with Quantitative and Adaptive Hot-Bytes Identification." That summary explicitly states that the paper’s actual system name is **Finch**, while referring to it as **cuFuzz/Finch** because of the query wording. Finch is a different system focused on **multi-objective optimization over branch distances**, **min-Pareto seed selection**, and **neural hot-byte prediction**, with experiments on **10 real-world programs** and **LAVA-M**. It should therefore not be conflated with either CUDA-oriented cuFuzz or the fork-awareness framework [2307.02289].

Taken together, the literature uses **cuFuzz** for materially different contributions: whole-program GPU fuzzing with host-device coverage fusion, compiler-led CUDA-to-CPU bug-preserving transformation for fuzzing, and fork-awareness evaluation for coverage-guided fuzzers. Any precise technical discussion must therefore identify the paper and architecture in view rather than relying on the name alone.

Source: https://www.emergentmind.com/topics/cufuzz-43bf24bb-0e51-4990-a245-3f3954ad9c4d