cuFuzz: Multi-Use Fuzzing Framework

• cuFuzz is a term applied to multiple fuzzing artifacts, including a CUDA-oriented whole-program fuzzer, a compiler-runtime framework for translating CUDA to CPU, and an evaluation framework for fork-awareness.
• The CUDA fuzzer integrates AFL++ with NVBit to unify host-device coverage, effectively reducing false positives and improving crash detection in complex GPU-CPU interactions.
• The transformation-based approach employs LLVM IR translation along with PREX and AXIPrune optimizations to accelerate performance while preserving critical memory bugs.

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 (Ziad et al., 12 Mar 2026, Singh et al., 3 Jan 2026, Maugeri et al., 2023).

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 (Ziad et al., 12 Mar 2026, Singh et al., 3 Jan 2026, Maugeri et al., 2023, Nguyen et al., 2023).

Usage	Technical description	Paper
cuFuzz	CUDA-oriented, whole-program, coverage-guided fuzzer	(Ziad et al., 12 Mar 2026)
CuFuzz	Compiler-runtime co-design framework transforming CUDA to CPU for fuzzing	(Singh et al., 3 Jan 2026)
cuFuzz	Fork-awareness evaluation idea for coverage-guided fuzzers	(Maugeri et al., 2023)
Finch	Quantitative and adaptive hot-bytes identification; actual system name is Finch	(Nguyen et al., 2023)

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 (Ziad et al., 12 Mar 2026).

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 (Ziad et al., 12 Mar 2026).

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 (Ziad et al., 12 Mar 2026).

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 (Singh et al., 3 Jan 2026).

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 (Singh et al., 3 Jan 2026).

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 (Maugeri et al., 2023).

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 (Nguyen et al., 2023).

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.