---
title: 'CudaForge: Dual-Use CUDA Framework'
url: https://www.emergentmind.com/topics/cudaforge
type: topic
---

# CudaForge: Dual-Use CUDA Framework

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 [2511.01884]. 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 [2109.00673]. 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” [2601.15727].

## 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 [2109.00673].

## 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 [2511.01884].

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 [2511.01884].

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 $S_m$ above the $75$th 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` [2511.01884].

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
$$
\text{tolerance} = 1\text{e}{-4}.
$$
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\(_1\), median speedup, and 75th percentile speedup [2511.01884].

## 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$\times$ average speedup**, **1.107 median speedup**, **1.592$\times$ 75th percentile speedup**, and **70.8% Fast\(_1\)**. By level, the reported outcomes are **96% correctness, 1.448$\times$ speedup** for Level 1; **100% correctness, 2.104$\times$ speedup** for Level 2; and **96% correctness, 1.283$\times$ speedup** for Level 3. On the stratified subset $\mathcal{D}^*$ of 25 tasks—10 from Level 1, 10 from Level 2, and 5 from Level 3—the framework reports **100% correctness**, **1.767$\times$ average speedup**, **1.322 median speedup**, **1.736$\times$ 75th percentile speedup**, and **84% Fast\(_1\)** [2511.01884].

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$\times$ performance** for OpenAI-o3, **90.8% correctness and 1.107$\times$ performance** for `o3-self-refine`, **97.6% correctness and 1.222$\times$ performance** for `o3-correction`, and **88.4% correctness and 1.509$\times$ 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$\times$ speedup**, whereas CudaForge is reported at **98.0% correctness and 1.662$\times$ 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 [2511.01884].

The framework also reports substantial **cross-GPU generalization**. On the same subset setting, it achieves **100% correctness and 1.767$\times$ performance** on RTX 6000, **100% correctness and 1.327$\times$ performance** on RTX 4090, **100% correctness and 1.841$\times$ performance** on A100, and **100% correctness and 1.320$\times$ performance** on RTX 3090. It further reports **cross-model generalization**, including **O3 / GPT-5: 100% correctness, 2.114$\times$ performance, 96% Fast\(_1\)**; **O3 / Claude-Sonnet-4: 100% correctness, 1.829$\times$ performance**; **O3 / GPT-OSS-120B: 100% correctness, 1.364$\times$ performance**; **GPT-5 / O3: 100% correctness, 1.896$\times$ 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 \$0.3 API cost**, contrasted with a prior agentic baseline reported at **6 H100 GPU hours** and **\$5 API cost per kernel** [2511.01884].

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$\times$ to 2.42$\times$**. In round 5, a correctness failure is traced to an uninitialized `target_logit` in thread 0, which is then broadcast using `_shfl_sync`. 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 [2511.01884].

## 4. CudaForge as a CUDA-to-RISC-V execution pipeline

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 [2109.00673].

The resulting pipeline has four stages:

1. **CUDA source code $\rightarrow$ NVVM IR**  
2. **NVVM IR $\rightarrow$ SPIR-V IR**  
3. **SPIR-V IR $\rightarrow$ OpenCL IR**  
4. **OpenCL IR $\rightarrow$ 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 [2109.00673].

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 [2109.00673].

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 [2109.00673].

## 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`** [2109.00673].

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 [2109.00673].

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”** [2601.15727].

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
$$
fast_p = \frac{1}{N} \sum_{i=1}^{N} \mathbbm{1}(\text{correct}_i \land \{\text{speedup}_i > p\}),
$$
and emphasizes functional passes before optimization passes [2511.13274]. 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 [2603.12440].

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 a synthesis system, but it is directly relevant to CudaForge-type optimization because it provides kernel fingerprints and similarity measures that can inform autotuners and compilers [1707.02423].

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**.

Source: https://www.emergentmind.com/topics/cudaforge