CudaForge: Dual-Use CUDA Framework

Updated 4 July 2026

CudaForge is a dual-use framework comprising a training-free multi-agent optimizer for CUDA kernels and an end-to-end CUDA-to-RISC-V translation pipeline.
The 2025 system employs a Coder-Judge loop that generates and refines kernels using hardware feedback (e.g., Nsight Compute metrics) to maximize speed and correctness.
The 2021 system translates CUDA code into intermediate representations (NVVM, SPIR-V, OpenCL) to enable execution on open-source RISC-V GPUs, ensuring legacy code portability.

CudaForge is the name of two distinct CUDA-related systems in the recent literature. In one usage, it denotes a training-free, multi-agent framework for CUDA kernel generation and optimization that separates synthesis and evaluation into a Coder-Judge loop and incorporates hardware feedback such as Nsight Compute metrics and GPU specifications (Zhang et al., 23 Oct 2025). In another, it denotes an end-to-end CUDA-to-RISC-V execution pipeline that translates CUDA source code through NVVM IR, SPIR-V IR, and OpenCL IR before execution on the open-source Vortex RISC-V GPGPU via POCL-based runtime support (Han et al., 2021). A later survey of automated kernel generation explicitly places CudaForge within the agentic optimization literature and characterizes it as a system that “employs a Coder-Judge loop driven by hardware-level feedback” (Yu et al., 22 Jan 2026).

1. Name, scope, and disambiguation

Within the literature considered here, “CudaForge” does not refer to a single unified artifact but to two systems that address different layers of the CUDA stack. The earlier system addresses portability and execution of existing CUDA programs on open-source hardware. The later system addresses automatic synthesis and optimization of CUDA kernels for performance-critical AI workloads.

CudaForge usage	Primary objective	Core mechanism
2021 system	Execute CUDA source code on an open-source RISC-V GPU architecture	CUDA source $\rightarrow$ NVVM IR $\rightarrow$ SPIR-V IR $\rightarrow$ OpenCL IR $\rightarrow$ RISC-V binary / execution on Vortex
2025 system	Generate and optimize efficient CUDA kernels	Training-free multi-agent Coder-Judge loop with Nsight Compute and GPU-spec feedback

This bifurcation is significant because the two systems solve orthogonal problems. The 2021 CudaForge seeks to preserve the CUDA programming model while decoupling execution from NVIDIA hardware and proprietary runtime dependence. The 2025 CudaForge seeks to automate the expert workflow of kernel engineering—initial implementation, correctness debugging, profiling, bottleneck analysis, and iterative optimization—without retraining the underlying model (Han et al., 2021).

2. CudaForge as a multi-agent CUDA kernel optimization framework

The 2025 CudaForge is defined as a training-free, multi-agent workflow for CUDA kernel generation and optimization. Its objective is to produce a kernel that is both functionally equivalent to the PyTorch reference and as fast as possible on the target GPU. The framework is explicitly modeled on the workflow of human CUDA experts: develop an initial kernel, test correctness, inspect profiling feedback, diagnose bottlenecks, and iteratively refine the code (Zhang et al., 23 Oct 2025).

Its architecture consists of two independent LLM agents. The Coder generates candidate CUDA kernels. The Judge evaluates those kernels and returns feedback. The Coder receives the task description, the PyTorch reference, the previous kernel candidate, and the Judge’s feedback; it then proposes a revised candidate. That candidate is compiled and executed on test cases, and, if correct, it is profiled with Nsight Compute. The Judge then operates in one of two modes. In correction mode, it reads error logs, mismatched outputs, and the current CUDA code, then identifies the single most critical correctness issue. In optimization mode, it uses the kernel, GPU specifications, and Nsight Compute metrics to identify the dominant bottleneck and propose one optimization method together with a concrete modification plan. The main setting uses a maximum of $N=10$ rounds (Zhang et al., 23 Oct 2025).

A central design choice is the use of hardware feedback rather than text-only self-refinement. The Judge can inspect GPU architecture and limits, including memory bandwidth, register limits, and shared memory capacity, together with Nsight Compute metrics such as memory throughput, occupancy, warp efficiency, scoreboard stalls, memory dependency stalls, barrier stalls, and branch resolving stalls. The purpose is to infer whether a kernel is memory-bound, compute-bound, occupancy-limited, register-pressure limited, or latency-bound. The paper argues that feeding the entire Nsight Compute metric space degrades optimization quality; accordingly, CudaForge filters profiling data to a 24-metric key subset using runtime correlation analysis. For metric $m$ , the global correlation score is defined as

$S_m = \frac{1}{n}\sum_{i=1}^{n} |r_{m,i}|,$

where $r_{m,i}$ is the Pearson correlation between metric $m$ and runtime on task $i$ . Metrics with $\rightarrow$ 0 above the $\rightarrow$ 1th percentile are retained, yielding 24 distinct metrics. Representative entries include sm__cycles_active.avg, launch__occupancy_limit_registers, dram__throughput.avg.pct_of_peak_sustained_elapsed, and smsp__warp_issue_stalled_long_scoreboard_per_warp_active.pct (Zhang et al., 23 Oct 2025).

Correctness is defined by a two-stage process: successful compilation followed by execution on predefined test cases, with outputs matching the PyTorch reference within tolerance

$\rightarrow$ 2

Performance is then evaluated as the speed ratio between the generated kernel and the PyTorch baseline on the same GPU. The reported metrics include correctness, performance, Fast $\rightarrow$ 3, median speedup, and 75th percentile speedup (Zhang et al., 23 Oct 2025).

3. Empirical results, generalization, and operating characteristics

The principal evaluation of the 2025 CudaForge is conducted on KernelBench, comprising 250 tasks: 100 in Level 1, 100 in Level 2, and 50 in Level 3. On the full KernelBench benchmark on RTX 6000, CudaForge reports 97.6% correctness, 1.677 $\rightarrow$ 4 average speedup, 1.107 median speedup, 1.592 $\rightarrow$ 5 75th percentile speedup, and 70.8% Fast $\rightarrow$ 6. By level, the reported outcomes are 96% correctness, 1.448 $\rightarrow$ 7 speedup for Level 1; 100% correctness, 2.104 $\rightarrow$ 8 speedup for Level 2; and 96% correctness, 1.283 $\rightarrow$ 9 speedup for Level 3. On the stratified subset $\rightarrow$ 0 of 25 tasks—10 from Level 1, 10 from Level 2, and 5 from Level 3—the framework reports 100% correctness, 1.767 $\rightarrow$ 1 average speedup, 1.322 median speedup, 1.736 $\rightarrow$ 2 75th percentile speedup, and 84% Fast $\rightarrow$ 3 (Zhang et al., 23 Oct 2025).

The baseline comparisons are intended to separate the effects of raw model capability, self-refinement, correctness feedback, and hardware-aware optimization. The paper reports 57.6% correctness and 0.680 $\rightarrow$ 4 performance for OpenAI-o3, 90.8% correctness and 1.107 $\rightarrow$ 5 performance for o3-self-refine, 97.6% correctness and 1.222 $\rightarrow$ 6 performance for o3-correction, and 88.4% correctness and 1.509 $\rightarrow$ 7 performance for o3-optimization. Against Kevin-32B on tasks comparable to KernelBench Levels 1–2, Kevin-32B is reported at 82.0% correctness and 1.10 $\rightarrow$ 8 speedup, whereas CudaForge is reported at 98.0% correctness and 1.662 $\rightarrow$ 9 performance. These comparisons are used to argue that correctness feedback alone stabilizes generation but does not maximize speed, while optimization feedback alone is insufficient if kernels remain functionally broken (Zhang et al., 23 Oct 2025).

The framework also reports substantial cross-GPU generalization. On the same subset setting, it achieves 100% correctness and 1.767 $\rightarrow$ 0 performance on RTX 6000, 100% correctness and 1.327 $\rightarrow$ 1 performance on RTX 4090, 100% correctness and 1.841 $\rightarrow$ 2 performance on A100, and 100% correctness and 1.320 $\rightarrow$ 3 performance on RTX 3090. It further reports cross-model generalization, including O3 / GPT-5: 100% correctness, 2.114 $\rightarrow$ 4 performance, 96% Fast $\rightarrow$ 5; O3 / Claude-Sonnet-4: 100% correctness, 1.829 $\rightarrow$ 6 performance; O3 / GPT-OSS-120B: 100% correctness, 1.364 $\rightarrow$ 7 performance; GPT-5 / O3: 100% correctness, 1.896 $\rightarrow$ 8 performance; Claude-Sonnet-4 / O3: 88% correctness; and QwQ / O3: 84% correctness. Operationally, the reported average cost is about 26.5 minutes per optimized kernel on one RTX6000 and about \$\rightarrow$95 API cost per kernel (Zhang et al., 23 Oct 2025).

A concrete case study on CrossEntropyLoss illustrates the framework’s feedback loop. In round 2, the Judge detects barrier stalls and recommends replacing block-level synchronization with warp-level shuffle reduction, after which speed improves from 1.66$N=10$0 to 2.42$N=10 $1. In round 5, a correctness failure is traced to an uninitialized <code>target_logit</code> in thread 0, which is then broadcast using <code>_shfl_sync</code>. In rounds 6–7, the Judge identifies long-scoreboard stalls due to global-memory latency and high register pressure, and recommends lowering register use and caching logits in shared memory (<a href="/papers/2511.01884" title="" rel="nofollow" data-turbo="false" class="assistant-link" x-data x-tooltip.raw="">Zhang et al., 23 Oct 2025</a>). <h2 class='paper-heading' id='cudaforge-as-a-cuda-to-risc-v-execution-pipeline'>4. CudaForge as a CUDA-to-RISC-V execution pipeline</h2> The 2021 CudaForge addresses a different problem: making existing CUDA source code executable on an open-source RISC-V GPU architecture. Its motivation is that CUDA is the dominant GPU programming model for scientific computing and machine learning, but it is tightly coupled to NVIDIA’s proprietary ecosystem, while RISC-V lacks equivalent software support. The explicit objective is to reuse existing CUDA code without rewriting it into another high-level language such as SYCL or OpenCL. The authors reject both source-to-source migration as insufficient and direct CUDA-to-RISC-V compilation as difficult to maintain as RISC-V evolves and accumulates extensions (<a href="/papers/2109.00673" title="" rel="nofollow" data-turbo="false" class="assistant-link" x-data x-tooltip.raw="">Han et al., 2021</a>). The resulting pipeline has four stages: <ol> <li>CUDA source code$ N=10$2 NVVM IR

NVVM IR $N=10$3 SPIR-V IR

SPIR-V IR $N=10$4 OpenCL IR

OpenCL IR $N=10$5 RISC-V binary / execution on Vortex

In the architecture described in the paper, the project-specific components are the NVVM-to-SPIR-V translator and RISC-V runtime/library support. The rationale for introducing SPIR-V as an intermediate layer is that it decouples the CUDA front end from the RISC-V GPU back end, improves maintainability, and reduces sensitivity to ongoing ISA evolution. The end-to-end flow compiles CUDA into NVVM IR using the CUDA toolkit, translates NVVM IR into SPIR-V, converts SPIR-V into OpenCL IR via the LLVM-SPIRV translator, feeds OpenCL IR into POCL, links it with a RISC-V backend runtime/library, and executes the resulting object on Vortex, the RISC-V GPGPU (Han et al., 2021).

The system treats CUDA as C++ plus GPU-specific extensions. Ordinary C++ is not itself the problem; CUDA-specific constructs such as __global__, built-in thread/block intrinsics like threadIdx.x, memory hierarchy assumptions, synchronization, and atomic operations are the critical elements. CUDA kernels are first lowered to NVVM IR, where !nvvm.annotations marks kernel entry points and functions such as llvm.nvvm.read.ptx.sreg.tid.x() retrieve thread identifiers. The translator then handles three content types in NVVM IR: metadata, NVVM built-in function declarations, and device-independent instructions. Backend-independent metadata are copied, CUDA-specific metadata are rewritten or dropped, nvvm.annotations is converted into SPIR-V-style kernel metadata, and kernel_arg_addr_space is added to record that kernel pointer arguments point to global memory. The built-in translation replaces NVVM calls such as llvm.nvvm.read.ptx.sreg.tid.x() with SPIR-V built-in-variable accesses followed by component extraction, using the dimension mapping x \rightarrow 0, y \rightarrow 1, z \rightarrow 2. Device-independent instructions are largely reused through the existing OpenCL-SPIR-V translator rather than reimplemented (Han et al., 2021).

The paper also describes feature support for multithreading, barrier synchronization, and atomic instructions. NVVM block and thread index built-ins are mapped to SPIR-V/OpenCL built-in variables such as BuiltInWorkgroupId and BuiltInLocalInvocationId, enabling multi-dimensional grid and block indexing. NVVM’s llvm.nvvm.barrier0() is translated into an OpenCL barrier call with CLK_LOCAL_MEM_FENCE|CLK_GLOBAL_MEM_FENCE to preserve CUDA-style block synchronization semantics. CUDA atomics such as atomicAdd, atomicSub, and atomicExch, represented in NVVM via atomicrmw, are translated into corresponding OpenCL atomic function calls because the argument prototypes are compatible (Han et al., 2021).

5. Evaluation and limitations of the CUDA-to-RISC-V CudaForge

Evaluation of the 2021 CudaForge uses the Rodinia benchmark suite and reports benchmark-level support status for different CUDA features.

Application	Feature	Support
b+tree	-	yes
bfs	-	yes
cfd	double3 type	yes
huffman	atomic	yes
pathfinder	memory hierarchy	yes
gaussian	-	yes
hotspot	-	yes
hotspot3D	-	yes
lud	memory hierarchy	yes
nw	-	yes
streamcluster	-	yes
particlefilter	d2i	on going
backprop	`__log2f`	on going
lavaMD	d2i	on going
kmeans	texture	no
hybrid sort	texture	no
leukocyte	texture	no

The reported pattern is that most benchmarks with ordinary control flow, multithreading, memory hierarchy handling, and atomics were supported, whereas unsupported or partially supported cases involve texture memory, certain mathematical functions, narrowing conversions such as d2i, and specific functions such as __log2f (Han et al., 2021).

The implementation is described as lightweight, built on LLVM, capable of running without NVIDIA GPUs, and dependent on existing open-source components including the CUDA toolkit, LLVM/SPIRV translator, POCL, and Vortex. The design trades direct compilation for modularity and portability: using intermediate representations adds translation stages and dependence on existing compiler infrastructure, but it improves maintainability and decouples front-end and back-end evolution. The paper identifies immediate future work as supporting the remaining Rodinia benchmarks, adding support for texture memory and mathematical functions, and extending the corresponding runtime libraries needed to execute those features on RISC-V GPUs (Han et al., 2021).

Taken in context, the significance of this earlier CudaForge is not that it provides a fully general CUDA compiler, but that it demonstrates a practical route for bringing CUDA workloads into the open-source hardware ecosystem while preserving the original CUDA programming model. This suggests a portability strategy centered on IR translation and runtime adaptation rather than source rewriting.

6. Position in the broader kernel-generation and CUDA-analysis literature

The survey “Towards Automated Kernel Generation in the Era of LLMs” treats CudaForge as part of the emerging literature on LLM-driven and agentic approaches for generating and optimizing GPU kernels. That survey organizes the space around supervised fine-tuning, reinforcement learning, agentic optimization workflows, datasets and knowledge bases, benchmarks, and evaluation metrics. Within its taxonomy of agentic systems, it lists learning mechanisms, external memory management, hardware profiling integration, and multi-agent orchestration, and explicitly states that “CudaForge employs a Coder-Judge loop driven by hardware-level feedback”. It also identifies the general optimization pattern as a “closed-loop, self-improving paradigm” and an “iterative, feedback-driven loop” (Yu et al., 22 Jan 2026).

In this broader context, CudaForge belongs to a family of systems that treat kernel generation as performance-oriented program synthesis rather than mere correctness-oriented code completion. KForge, for example, is a platform-agnostic two-agent system with a generation agent and a performance analysis agent that synthesizes programs across NVIDIA CUDA and Apple Metal, formalizes its correctness-and-speed metric as

$N=10$ 6

and emphasizes functional passes before optimization passes (Sereda et al., 17 Nov 2025). KernelFoundry, by contrast, treats GPU kernel optimization as an evolutionary search problem, combining MAP-Elites quality-diversity search, meta-prompt evolution, and template-based parameter optimization, and explicitly lists CudaForge as a prior multi-agent CUDA kernel optimizer with “coder + judge” and hardware feedback (Wiedemann et al., 12 Mar 2026).

A complementary but distinct line of work is represented by CUDAflow, which performs static binary analysis, dynamic PC sampling, CFG construction, and subgraph matching to characterize CUDA kernels by control-flow structure, transition behavior, instruction mix, and register allocation. CUDAflow is not an overview system, but it is directly relevant to CudaForge-type optimization because it provides kernel fingerprints and similarity measures that can inform autotuners and compilers (Lim et al., 2017).

These connections indicate that the newer CudaForge occupies a specific niche within automated kernel engineering: it is training-free, multi-agent, profiling-aware, and CUDA-centric, whereas adjacent systems either broaden the backend space, introduce evolutionary search, or focus on kernel characterization rather than generation. The older CudaForge occupies a different niche altogether, centered on CUDA portability onto RISC-V GPUs through IR-level translation. The shared name therefore spans two historically and technically distinct strands of CUDA systems research: agentic optimization and cross-ISA execution support.