Litespark-Nemotron-3.5 Super VL: Frontier CUDA Kernel Generation from PyTorch Code

Motivation

Language models now write CUDA routinely, and much of what they write depends on other factors such as vendor libraries that handle matrix multiplication, launch configurations that have been hard-coded for the hardware and in some cases an instruction set that is available on one architecture but quietly appears in the inner loop. Each of these makes generation easier but also results in narrower implementation.

This work seeks to address a prevailing problem in modern performance engineering. Can a model read a PyTorch module and write a raw CUDA kernel that computes the same forward pass within a fixed numerical tolerance while remaining architecture agnostic?

A kernel with no dependencies and no assumptions about the hardware is a portable starting point. The aim is to translate from PyTorch once, and get the result to run correctly on any modern NVIDIA GPU. Making it fast on a particular hardware is a separate problem which requires several code iterations and hence is out of scope of this project. On the other hand, accurately translating PyTorch code to a kernel that compiles with numerical correctness is a fundamental step in AI performance engineering.

The model performing the translation is open-weights, adapted from NVIDIA Nemotron 3.5 Super VL open model with parameter-efficient fine-tuning. If the capability exists, it therefore sits in weights that can be run, inspected and trained further in-house, rather than behind an API. The natural comparison is against externally hosted models given the same task.

Methodology

The model receives the source of a PyTorch module and returns CUDA device code that computes the module's forward pass. Four requirements govern what that code may do:

• No PyTorch: The kernel operates on plain device memory. It cannot call back into the framework it is replacing, so every part of the computation has to be expressed in CUDA itself.
• No vendor libraries: A convolution or a matrix multiplication is written out in full, which means the model cannot delegate the hardest part of the problem; it has to understand the operation well enough to implement it.
• No architecture-specific features: Nothing in the kernel may depend on an instruction, a cache size or a core count belonging to one GPU generation.
• No assumptions about launch configuration: The kernel never decides how it is launched. The configuration is set outside the kernel at runtime for whichever GPU the kernel runs on, so the kernel must produce the right answer under whatever it is given. The standard way to achieve this is a grid-stride loop:

__global__ void scale(const float* x, float* y, float a, int n) {
    for (int i = blockIdx.x * blockDim.x + threadIdx.x;
         i < n;
         i += gridDim.x * blockDim.x)
        y[i] = a * x[i];
}

Each thread starts at its own index and advances by the total number of threads launched, so every element is covered exactly once whether the launch provides one thread per element or far fewer. The tempting alternative, computing the index as blockIdx.x * 256 + threadIdx.x, is correct only when the block happens to contain 256 threads, and wrong without any error being raised whenever it does not. The kernels are expected to be straightforward and are judged on correctness alone.

The model

The translator is NVIDIA Nemotron 3.5 Super VL, the new 120-billion-parameter model of which only 12 billion work on any given token. It is a mixture of experts, so most of the models sit out each step and wait patiently for their turn. It is a multimodal model that can read text, images and video. For this job we only needed it to read Python and write CUDA, but clearly the model is capable of much more.

We adapted it with parameter-efficient fine-tuning on paired examples: a PyTorch module and a hardware-agnostic CUDA kernel that computes the same forward pass. Generation runs with reasoning disabled. The model reads the module and writes the kernel in a single pass, with no intermediate drafts.

How the generated kernels are tested

As many generations of coders know very well, compilation alone proves little, since incorrect code compiles as good as correct code. Each kernel must therefore pass six checks, and each check exists because a weaker version would admit an incorrect kernel.

The reference is computed at higher precision. Floating-point addition is not associative: summing the same values in a different order yields a slightly different result. A kernel that does so is not wrong, so exact agreement with PyTorch is the wrong standard. Correctness instead means agreement with a fixed target within a set tolerance, and that target must carry less rounding error than the kernel it judges. We compute it one precision level higher, so FP32 kernels are graded against FP64 results. Inputs and reference outputs are computed once and frozen, so every kernel for a given problem is graded on identical data.

Tolerances are fixed before grading. Element-wise operations receive the tightest tolerance, operations that accumulate many values receive more, and numerically sensitive operations such as exponentials receive more still. Each operation takes its tolerance from a fixed table written before any kernel is graded; an operation absent from the table receives the strictest value. The two errors are not symmetric. A tolerance set too tight fails a correct kernel visibly; one set too loose passes an incorrect kernel silently.

The worst element decides. We measure error at every output element and judge the kernel by the largest. A single incorrect value, such as one produced by a boundary error at the edge of a tensor, fails the kernel regardless of how many others are correct.

The worst input decides. Each kernel runs on several independently seeded inputs, and each seed draws new weights as well as new data, so the kernel faces distinct problems rather than one problem with different values. Where an operation behaves differently for negative values, inputs include them, so a kernel that handles only one branch of a conditional is caught. The kernel's score is its lowest across seeds. One run is repeated to confirm the output does not vary between identical launches.

Both architectures must pass. Every kernel is compiled and graded on Ada and Ampere GPUs, and keeps the lower of the two scores. A kernel that relies on an instruction one architecture lacks fails on that architecture. This converts the portability requirement from a stated rule into a tested one.

The kernel must perform the computation. Source that calls PyTorch is rejected before execution. Kernels that copy an input to the output, or return the same result for inputs whose correct outputs differ, are also rejected. The entire attention has to be on the kernel itself.

Evaluation

Problems are drawn from the first two levels of KernelBench (https://kernelbench.com/). Level 1 problems are single operations: a convolution, a matrix multiplication, a reduction. Level 2 problems chain several operations into one module (a convolution followed by a bias, an activation and a normalization, for example), and the kernel must implement the whole chain. Each model makes one attempt per problem: no retries, and no selection among multiple samples.

We compare three groups: the post-trained model, the base model, externally hosted frontier and open-weights models. Reasoning is disabled for the adapted model to focus on the role of the weights and what the model actually learnt instead of test-time search. For the others, it is disabled where the provider permits and set to its minimum where it does not, keeping conditions as close to equal as each provider allows.

The post-trained model receives the prompt it was trained on; all other models, including the base model, receive a common prompt stating the same four requirements. The difference between the model before and after adaptation therefore reflects fine-tuning together with its prompt, not fine-tuning alone.

A kernel passes only if its output is within tolerance at every element, on every seed, on both architectures. A small number of problems could not be graded for reasons unrelated to the submitted kernel and are excluded. Pass rates are reported over the remaining problems with 95% confidence intervals.

Results

In the figures, the post-trained model appears as Litespark-Nemotron-3.5-Super VL, and the base model as Nemotron-3.5-Super VL.

Level 1:

On single operations, the post-trained model passes 83.9% of problems, up from 72.7% before adaptation: a gain of 11.2 percentage points from fine-tuning.

The top five models span 79.5% to 83.9%, and at this problem count, those differences are not statistically separable. Level 1 establishes that an open-weights model, adapted with parameter-efficient fine-tuning, writes library-free CUDA for single operations as reliably as the strongest models available.

Level 2:

In the Level 2 evaluation, the post-trained model passes 82.1% of problems; the strongest hosted model, GPT-5.6 Sol, passes 57.0%. Before post-training, the same model passed 22.8%, so post-training lifts it by 59.3 points on composed problems, against 11.2 on single ones.

Set side by side, the two figures show how much harder composition is. GPT-5.6 Sol falls from 82.6% on single operations to 57.0% on composed ones, Kimi K3 from 82.3% to 42.5%, and GLM 5.2 from 69.9% to 20.2%. The adapted model moves from 83.9% to 82.1%.

Discussion

Adaptation is what helps solve the hardest problem

Level 1 asks whether a model can implement an operation in raw CUDA. Level 2 asks whether it can implement several in one kernel, carrying each intermediate result correctly into the next stage with no library to delegate to. Most capable models handle the first well. On the second, the hosted models decline sharply while the post-trained model holds steady. This suggests fine-tuning taught the model not only to write individual operations but to compose them.

The result comes from a single attempt without reasoning.

Every figure above reflects one generation per problem, and as we mentioned reasoning is either disabled or minimized. The post-trained model reaches its result in one pass, from weights that can be run and further trained in-house, that allows to reduce cost especially in situations where thousands of kernels are called at the same time.

Composed kernels are the input the next stage requires. Most of the speedup available over torch.compile comes from fusing operations the compiler leaves separate, which is the structure of a Level 2 problem. A model that writes correct composed kernels at this rate is the right starting point for follow up optimizations.

Limitations

While we are proud of what we have achieved and showed so far, this work has two main limitations. First, it measures correctness, not speed: a kernel passes if it computes the right answer, however slowly it runs, and performance is the subject of the next part of this work. Second, portability is established only on Ada and Ampere, the two architectures every kernel was graded on. We think that other NVIDIA architectures will behave the same way, since the same requirements apply, and we plan to extend the evaluation to confirm this.

Future work

The kernels produced here are correct and portable by construction, which makes them a starting point rather than an endpoint. The next part of this work takes them, together with the original PyTorch module, and optimizes for a named target GPU against torch.compile, keeping correctness and performance as separate problems, solved in sequence.

What is better than being able to do something difficult reliably? It is to do it FAST. The kernels produced here are correct and portable, which makes them the ideal starting point for the next part of this work that takes them, together with the original PyTorch module, and optimizes them for a named target GPU against torch.compile, treating correctness and performance as separate problems.

Conclusion

This work asked whether a model can read a PyTorch module and write a raw CUDA kernel that computes the same forward pass, with no vendor libraries, no architecture-specific features and no assumptions about launch configuration. Under strict grading, the answer is a strong yes. An open-weights model, adapted with parameter-efficient fine-tuning, matches the strongest hosted models on single operations and leads them by a wide margin on composed ones, passing 82.1% of Level 2 problems against GPT 5.6 Sol at 57.0%. It does so in a single pass, without reasoning, from weights that can be run and trained further in-house. The model weights are available on Hugging Face.