CASS: Nvidia to AMD Transpilation with Data, Models, and Benchmark

🎯 Pitch

General-purpose LLMs like GPT-4o and Claude score 0% on translating Nvidia GPU assembly to AMD—they completely lack the capability. But domain-specific models trained on CASS's 70k aligned code pairs achieve up to 37.5% assembly accuracy and match native performance in over 85% of source-level translations, while assembly success jumps from 0% to 100% in certain domains like physics.

1. Executive Summary

This paper introduces CASS, the first large-scale dataset and model suite for cross-architecture GPU code transpilation, targeting both source-level (CUDA ↔ HIP) and assembly-level (Nvidia SASS ↔ AMD RDNA3) translation across 70k verified, semantically aligned code pairs spanning 16 GPU domains. The CASS model family — domain-specific LLMs fine-tuned from Qwen2.5-Coder at 1.5B, 3B, and 7B scales — achieves 95% source translation accuracy and 37.5% assembly translation accuracy on the newly introduced CASS-Bench evaluation suite, substantially outperforming commercial baselines including GPT-4o, Claude (both at 0% assembly accuracy), and the static translator Hipify (by 7.5% at the source level), while preserving runtime and memory behavior within ±5.6% for over 85% of test cases. The CASS dataset pipeline — combining repository scraping, synthetic generation via persona-augmented LLM prompting, Hipify-based transpilation, and dual-stack compilation with host/device assembly disentanglement — establishes that domain-specific fine-tuning on aligned cross-vendor assembly pairs can impart low-level ISA translation capabilities entirely absent from general-purpose models, but only for domains well-represented in training, with assembly accuracy varying from 0% in math, data structures, and graph tasks up to 100% in physics simulations.

2. Context and Motivation

The Core Problem: GPU Vendor Lock-In and the Missing Infrastructure for Cross-Architecture Translation

The paper addresses a problem with two layers: a surface-level translation gap (we have CUDA code and want to run it on AMD GPUs) and a deeper infrastructure gap (we have no training data, no aligned assembly pairs, no benchmarks, and no models purpose-built for learning GPU-to-GPU translation). The combination means that even if we wanted to use machine learning to automate GPU transpilation — which is an inherently attractive idea given the complexity of assembly languages — there has been no foundation on which to build such a system.

The surface-level problem is what the authors term vendor lock-in in GPU computing. NVIDIA's CUDA has become the de facto programming model for GPU acceleration, but the ecosystem is vertically integrated: CUDA code compiles through NVIDIA's proprietary toolchain (nvcc → PTX → SASS) and executes only on NVIDIA hardware, because different GPU vendors implement incompatible instruction set architectures. This means organizations with large CUDA codebases — spanning scientific computing, machine learning, graphics, and high-performance computing — face two unattractive choices when considering AMD hardware: rewrite their entire codebase from scratch, or accept that their software is permanently tied to a single hardware vendor.

The practical stakes are significant. AMD GPUs offer potentially favorable performance-per-dollar economics (the paper cites benchmarks from AMD and The Verge, references [3] and [4]), and are increasingly adopted in both data centers and consumer devices (Section 1, citing reference [5]). If organizations could easily migrate their CUDA investments to AMD hardware, they would gain negotiating leverage and deployment flexibility. But the cost of manual migration — rewriting kernels, debugging correctness, and re-tuning performance — is prohibitive for large codebases, making the lock-in effectively absolute.

Why This Problem Matters: Beyond Source Translation

The paper distinguishes between two levels of translation that motivate different use cases, and this distinction is crucial for understanding why the problem matters:

Source-level translation (CUDA → HIP) is what existing tools attempt. AMD's HIP programming model is intentionally designed as a near-drop-in replacement for CUDA, with analogous API calls, memory management, and kernel launch semantics. HIPIFY, AMD's own static translator, converts CUDA source to HIP through syntactic substitution — replacing cudaMalloc with hipMalloc, cudaMemcpy with hipMemcpy, and so on. This enables recompilation of legacy CUDA code on AMD hardware, preserving the ability to make future modifications. It is valuable for codebases that are actively maintained.

Assembly-level translation (SASS → RDNA3) addresses a fundamentally different need, one that the paper argues is critical for democratizing the GPU landscape but has been entirely neglected. Assembly translation would enable:

Execution of precompiled CUDA binaries without source access. This matters for legacy software where the original source is lost, proprietary binaries distributed by third parties, or any situation where recompilation is impossible. Only assembly-level translation can address these cases.
Transfer of low-level, hardware-specific optimizations. CUDA developers often tune kernels at the assembly level to exploit NVIDIA-specific features — register allocation strategies, instruction scheduling, memory access patterns optimized for the memory hierarchy of specific NVIDIA architectures. These optimizations exist below the PTX intermediate representation and are invisible to IR-level tools. If we translate only at the source or IR level, we lose these optimizations; the AMD compiler must rediscover equivalent performance from scratch. Assembly-to-assembly translation, in principle, could map these optimizations directly onto AMD's instruction set — NVIDIA's register-blocking strategy becomes AMD's equivalent, NVIDIA's shared memory access pattern becomes AMD's LDS access pattern. The paper explicitly positions assembly-level translation as a way to "leverage hardware-specific optimizations below the intermediate representation (IR) level, that may be missing altogether in the corresponding AMD codebase" (Section 2.1).
Enabling automation beyond source-level rewrites. The paper frames assembly translation as part of a broader vision of cross-vendor hardware compatibility — analogous to how CPU binary translation (e.g., Apple's Rosetta for x86 → ARM) enabled architecture transitions without requiring application rewrites. GPU binary translation is substantially harder because GPU ISAs are less standardized and more divergent, but the goal is the same.

The tension between these two levels of translation creates a rich technical problem. Source translation has seen some tooling investment (HIPIFY, CuPBoP-AMD, ZLUDA) but remains incomplete and error-prone. Assembly translation has been completely unexplored because, as the paper argues, the necessary infrastructure simply did not exist: no paired, aligned dataset of equivalent GPU instructions across vendors, no benchmark to measure progress, and no training corpus on which to build machine learning approaches.

Prior Approaches and Their Limitations

The paper identifies four categories of prior work, each with specific shortcomings that CASS is designed to address:

1. Static source translators (HIPIFY). HIPIFY is AMD's official tool for converting CUDA source to HIP. It works by pattern-matching CUDA API calls and replacing them with HIP equivalents — essentially a sophisticated find-and-replace engine with awareness of GPU programming idioms. The paper reports two critical limitations. First, HIPIFY exhibits a high failure rate: the authors found that approximately 43.9% of CUDA files in their pipeline failed HIPIFY conversion (Section 3.3), a figure consistent with prior work by Zahid et al. (2024, reference [10]) that documented high failure rates in CUDA-to-HIP conversion. Failures arise from HIPIFY's inability to handle unsupported CUDA features, edge cases in template instantiation, or complex dependency chains. Second, HIPIFY operates only at the source level — it cannot touch precompiled binaries, which means it addresses only half the vendor lock-in problem. If your CUDA binary was compiled years ago and the source is unavailable, HIPIFY offers no help.

2. Intermediate representation translators (CuPBoP-AMD, ZLUDA). CuPBoP-AMD (Chen et al., 2023, reference [15]) translates NVIDIA's NVVM IR (LLVM-based) to HIP-compatible LLVM IR, operating one level below source in the compilation stack. ZLUDA (Janik, 2024, reference [6]) goes further, implementing a runtime system that intercepts CUDA API calls and dynamically translates embedded PTX or SASS into AMD-compatible code via LLVM. ZLUDA originally targeted Intel GPUs and now supports AMD RDNA3 through runtime patching. While ZLUDA is the closest prior work to assembly-level translation, the paper identifies a fundamental limitation: ZLUDA operates at the LLVM IR level, not the hardware assembly level. It translates PTX/SASS to LLVM IR, then lets AMD's backend compile that IR to RDNA3. This means ZLUDA cannot capture or exploit NVIDIA-specific optimizations that were applied below the IR level — the backend optimizations that the NVIDIA compiler performed when lowering PTX to SASS. The paper's framing is nuanced here: ZLUDA is "a reasonable level in the stack to target" (Section 2.1), but it leaves performance on the table because it only sees the intermediate representation, not the final, optimized hardware instructions. Assembly-to-assembly translation, by contrast, would work directly with the final optimized form on both sides.

The paper also notes that GPU Ocelot (Diamos et al., 2009, reference [18]) explored dynamic binary translation from CUDA to AMD/x86 ISAs at runtime, but "was limited by poor scalability and high overhead, making it impractical for modern GPU workloads." The authors cite Ocelot's lack of consistent updates as a additional failure mode shared across these tools: "All these tools have lacked consistent updates to keep up with CUDA advances, suffer from usability issues, and operate only at the source level" (Section 2.1).

3. CPU assembly transpilation with language models (CRT, Guess & Sketch). In the CPU domain, recent work has demonstrated that language models can learn to translate between ISAs given sufficient aligned training data. CRT (Heakl et al., 2024, reference [24]) is a lightweight transpiler from x86 assembly (CISC) to ARM (RISC), trained on paired sequences of instructions. Guess & Sketch (Lee et al., 2023, reference [25]) integrates language models with symbolic reasoning to translate between ARMv8 and RISC-V, achieving correctness through a generate-and-verify pipeline. These successes are directly motivating for CASS — they show that assembly-to-assembly translation is tractable with modern sequence models — but the paper identifies a critical gap that prevents their approach from applying to GPUs: the lack of a training dataset. The CPU work succeeded because large paired corpora of equivalent x86/ARM/RISC-V functions could be constructed from existing compiler output. No such dataset existed for GPU ISAs. As the paper states: "Given the lack of such a rich dataset in the GPU space, a primary goal of this work is to enable such an exploration and transpilation across GPU vendors" (Section 2.2).

4. Existing GPU datasets and benchmarks (ComputeEval, Rodinia, SHOC, etc.). Table 1 in the paper provides a systematic comparison of existing GPU datasets and benchmarks along characteristics that matter for translation. Every existing resource — ComputeEval (NVIDIA-focused, reference [19]), Rodinia (CUDA/OpenCL/OpenMP, reference [20]), SHOC (CUDA/OpenCL, reference [21]), PolyBench (CUDA/OpenCL, reference [22]), BabelStream (HIP/CUDA/OpenCL, reference [23]), and The Stack (nearly 200k CUDA files but no AMD coverage, reference [11]) — fails on one or more critical dimensions. None provide SASS assembly. None provide RDNA3 assembly. None provide paired, aligned data across NVIDIA and AMD codebases. The Stack contains a large number of CUDA files but no corresponding HIP or assembly equivalents, meaning it cannot be used to train a translation model that maps from one to the other.

The paper is explicit about the structural gap: "To the best of our knowledge, no existing dataset provides paired source- and assembly-level Nvidia-AMD code, hindering effective training and benchmarking" (end of Section 2.3). This is not merely an inconvenience — it means the entire field of GPU assembly translation has been blocked by a data availability problem, not a fundamental algorithmic limitation.

How CASS Positions Itself

The paper positions CASS as infrastructure for a new research direction, not as a finished solution to GPU transpilation. This framing is important for understanding the significance of the contributions. The 37.5% assembly accuracy is presented not as a production-ready result but as evidence that the dataset enables learning — and the 0% accuracy of all commercial baselines (GPT-4o, Claude, Gemini 2.0 Flash) on the same task is evidence of how far general-purpose models are from this capability without domain-specific training data.

The paper's positioning has four key elements:

First, it fills the dataset gap. CASS is explicitly "the first dataset purpose-built for cross-vendor GPU assembly translation" (Section 2.3). The 70k aligned pairs — covering CUDA/HIP source and SASS/RDNA3 assembly — provide the critical resource that prior CPU assembly translation work relied on but that did not exist for GPUs. Without this, the machine learning approach demonstrated by CRT and Guess & Sketch simply cannot be applied to GPUs.

Second, it introduces the benchmark gap. CASS-Bench provides 40 curated tasks across 16 GPU domains with execution-verified ground truth, giving the community a standardized way to measure progress. Prior benchmarks (Rodinia, SHOC, PolyBench) focused on runtime performance comparisons between hardware platforms, not on measuring the correctness of generated translations. CASS-Bench is designed specifically to evaluate whether a translated program produces the same output as its source — a fundamentally different evaluation criterion.

Third, it demonstrates feasibility through model training. By fine-tuning Qwen2.5-Coder on the CASS dataset and achieving non-trivial assembly translation accuracy where all baselines score 0%, the paper provides a proof of concept that the dataset imparts genuine cross-architecture translation capability. The ablation study (Table 4) further demonstrates that the synthetic data and OpenCL pipelines each contribute meaningfully to performance, validating the composability of the data collection strategy.

Fourth, it establishes the performance-fidelity connection. By measuring not just compilation success but also runtime and memory behavior of translated code (Section 6, Figure 11), the paper addresses a concern that purely syntax-focused translation metrics might miss: does the translated code actually run correctly and efficiently? The finding that over 85% of translated samples fall within ±5.6% of native performance for both memory and runtime establishes that correctness and efficiency are correlated for the trained model — it is not just producing plausible-looking assembly that happens to compile.

The paper explicitly acknowledges that this is foundational work, not a deployment-ready system. The limitations section (Section 7) identifies specific shortcomings — limited domain coverage, single host/device GPU pair per vendor, and performance inadequate for production — that frame CASS as enabling future work rather than solving the problem outright. This candor strengthens the paper's positioning: the contribution is the infrastructure (dataset, benchmark, models, pipeline) that makes systematic research on GPU transpilation possible for the first time.

The Underlying Assumption Worth Examining

The paper's entire approach rests on an assumption that is worth surfacing explicitly: that aligned assembly pairs carry sufficient semantic correspondence that a sequence-to-sequence model can learn the mapping. This is not obvious. SASS and RDNA3 are different ISAs with different register files, different memory models, different instruction granularity, and different optimization philosophies. A single LDG.E (NVIDIA load from global memory) might correspond to multiple RDNA3 instructions depending on addressing mode, coalescing behavior, and cache hierarchy. The paper's analysis in Section 4.1 — showing that AMD device assembly is on average twice as long as NVIDIA's, and that CHRF syntactic similarity scores are low — quantifies the extent of the divergence. The fact that training on the paired data produces any non-zero assembly accuracy at all (37.5% compared to 0% for all baselines) is non-trivial evidence that the mapping is learnable despite the divergence, which is itself a contribution of the work — it establishes that GPU ISAs are not so fundamentally incompatible that machine translation is hopeless.

3. Technical Approach

3.1 Reader Orientation

This paper constructs a three-part infrastructure system — a dataset pipeline, a benchmark suite, and a family of fine-tuned language models — that together enable the first machine learning approach to cross-architecture GPU code translation at both the source level (CUDA ↔ HIP) and assembly level (Nvidia SASS ↔ AMD RDNA3). The core problem it solves is the complete absence of aligned, paired training data for GPU-to-GPU transpilation, which has prevented the application of sequence-to-sequence learning to the vendor lock-in problem despite demonstrated success in CPU assembly translation (CRT, Guess & Sketch). The solution takes the shape of a data-first approach: build a scalable pipeline to generate aligned code pairs across the two GPU stacks, verify their semantic equivalence through compilation and execution, train models on the resulting corpus, and measure progress with a purpose-built benchmark — establishing that domain-specific fine-tuning on this paired data can impart low-level ISA translation knowledge entirely absent from general-purpose models.

3.2 Big-Picture Architecture (Diagram in Words)

The CASS system comprises five major components connected in a sequential pipeline, with the final output being both a training dataset and trained models:

Data Ingestion Layer — Gathers CUDA source code from two channels: scraping public repositories (top-200 repos from The Stack v2, yielding 24k usable samples after filtering) and synthetic generation using a persona-augmented LLM (Qwen2.5-Coder 32B) prompted with parameterized templates across 9 domain categories, yielding 46.3k compilable samples from 85k generated. An additional 6k OpenCL samples are collected from The Stack. This component solves the cold-start problem: there is no pre-existing corpus of aligned GPU assembly.
Transpilation and Compilation Engine — Converts the gathered CUDA source to HIP using AMD's HIPIFY tool, discarding the ~43.9% that fail conversion, then compiles both the original CUDA and the converted HIP through their respective toolchains (Nvidia: nvcc → PTX → SASS, extracted via cuobjdump; AMD: hipcc → RDNA3 assembly, with modified host/device separation) to produce paired, aligned source and assembly files. This component solves the alignment problem: it ensures each CUDA sample has a semantically equivalent HIP counterpart, down to the assembly level, because both are compiled from the same algorithm specification.
Filtering and Deduplication Layer — Retains only samples that compile successfully on both Nvidia and AMD pipelines (handling asymmetric failures where one stack succeeds and the other fails), deduplicates across the combined corpus, and applies length constraints (files with fewer than 10 lines or more than 7,000 lines are removed, as are files without CUDA kernel definitions or trivial boilerplate like "Hello World"). This produces the final 70,694 samples of CASS-Instruct.
CASS-Bench Curation Pipeline — A separate, manual verification pipeline that uses Claude-3.7 to generate CUDA implementations for 40 prompts across 16 GPU-centric domains, compiles and executes them on Nvidia hardware to obtain reference outputs, generates corresponding AMD code, and iteratively corrects mismatches until output equivalence is manually verified — then runs both through the compilation pipeline to extract aligned host and device assembly pairs with execution-verified ground truth.
Model Training and Inference Stack — Fine-tunes Qwen2.5-Coder at 1.5B, 3B, and 7B parameter scales on the CASS-Instruct dataset using LLaMA-Factory with DeepSpeed, Liger Kernel, and Paged AdamW optimizer, employing a 16K-token context window at training and RoPE extrapolation to 32.7K tokens at inference, producing two model variants (source translation and assembly translation) and evaluating them against CASS-Bench.

Information flows linearly through the first three components to produce the dataset, which then feeds the fourth component (benchmark creation) and the fifth component (training). At inference time, the trained model takes a CUDA source or SASS assembly as input and outputs HIP source or RDNA3 assembly, which is then compiled and executed on AMD hardware for verification.

3.3 Roadmap for the Deep Dive

First, the CUDA data ingestion layer (scraping and synthetic generation): The foundation of the entire pipeline — without diverse, compilable CUDA source, there is nothing to transpile. I will explain the repository scraping strategy (why whole-repository downloads matter), the synthetic generation strategy (what persona-augmented prompting is and why it was needed to supplement scraped data), and the filtering criteria that determine which samples survive to the next stage.
Second, the transpilation and compilation engine: This is the technical core of the paper. I will explain the HIPIFY conversion step and its failure modes, then detail the Nvidia and AMD compilation pipelines — which are architecturally different (Nvidia's is opaque, AMD's is transparent) — and the engineering required to extract and separate host and device assembly from each. The modified AMD pipeline (deferring device binary insertion) is a key design choice that enables independent translation of host and device code.
Third, the OpenCL pipeline: A parallel data source that bypasses the CUDA/HIP framework entirely, providing complementary coverage. This component demonstrates the generality of the approach — aligned assembly pairs can be generated without going through HIPIFY, using OpenCL as a unified source that compiles to both backends.
Fourth, the CASS-Bench curation process: Since the main dataset is generated automatically and may contain subtle errors that survive compilation, CASS-Bench provides a smaller but manually verified evaluation set. I will explain the iterative generation-and-verification protocol that ensures execution equivalence between Nvidia and AMD outputs.
Fifth, model training and inference: The supervised fine-tuning setup on Qwen2.5-Coder, including the specific hyperparameter choices (the aggressive learning rate of $1 \times 10^{-5}$ , the 16K context window, the gradient accumulation strategy), the assembly preprocessing (whitespace normalization for CUDA but not HIP), and the RoPE extrapolation that extends inference to 32.7K tokens.
Sixth, the critical numbers and design decisions bundled throughout: I will surface the specific hyperparameters, thresholds, and engineering choices that make each component work — the 43.9% HIPIFY failure rate, the 49.1% synthetic compilation success rate, the 23.7% improvement from repository-level scraping, the token reduction from -Os compilation, and the two-fold cross-validation and other evaluation details.

3.4 Detailed, Sentence-Based Technical Breakdown

This is primarily a data infrastructure and model training paper whose core idea is that aligned cross-vendor GPU assembly pairs — when collected at scale through a carefully engineered compilation pipeline — can serve as training data for sequence-to-sequence language models that learn to translate between GPU instruction set architectures, a capability that general-purpose models completely lack.

CUDA Code Scraping from Public Repositories

The first data source is public CUDA code scraped from The Stack v2 dataset (reference [27]), a large corpus of deduplicated, license-compliant source code extracted from public repositories. The authors go significantly beyond simple file collection: they use the dataset's metadata to identify the top 200 repositories with the highest number of CUDA files and download these repositories in their entirety, preserving the original directory structure and relative imports (Section 3.1).

Why repository-level download matters. GPU code rarely lives in isolation — a .cu file typically depends on header files defining kernel launch configurations, custom data structures, utility functions, and build configurations. If you scrape individual .cu files without their surrounding repository context, these dependencies break, and the code becomes uncompilable. The paper quantifies this effect: repository-level download "improved compilation success by 23.7% compared to isolated file scraping" (Section 3.1). This is a practical finding — the difference between a dataset that mostly fails to compile (and therefore cannot produce assembly) and one where a substantial fraction survives through the full pipeline.

Post-extraction filtering. After downloading the repositories, the authors apply a series of filters to the extracted CUDA files:

Length constraints: Files with more than 7,000 lines are removed (overly long files likely contain non-kernel code, are difficult to compile, and produce assembly that exceeds context windows). Files with fewer than 10 lines are removed (trivially short files lack meaningful kernel definitions and are not representative of real GPU workloads). These thresholds are practical engineering choices, not theoretically motivated — 7,000 lines is approximately the limit where file size starts causing downstream compilation timeouts and assembly extraction failures.
Boilerplate removal: "Naive boilerplate samples (e.g., 'Hello World')" are filtered out. These are files that contain minimal CUDA API calls but no actual parallel computation — they compile but produce trivial assembly that teaches the model nothing about GPU-to-GPU translation.
Kernel definition requirement: Files lacking CUDA kernel definitions (marked by the __global__ qualifier) are removed. This is the critical filter: without a kernel definition, the file is not GPU code in the sense that matters. It might contain only host-side CUDA API calls (memory allocation, data transfer) with the actual computation happening elsewhere — such files would produce host assembly but no device assembly, failing to generate the paired device assembly that is the primary value of CASS.

The final yield from this scraping pipeline is 24k usable CUDA samples (Section 3.1).

Synthetic Data Generation via Persona-Augmented LLM Prompting

Scraping alone produces an insufficiently diverse corpus — public CUDA repositories skew toward certain domains (cryptography, linear algebra) and underrepresent others (physics simulation, graph algorithms, sparse computation). To fill these gaps and increase overall dataset size, the authors synthesize additional CUDA code using a large language model with a structured prompt-generation strategy they call variable-augmented persona strategy (Section 3.2).

The prompting strategy. The authors define a set of natural language prompt templates, each containing variable placeholders enclosed in curly braces. These templates cover nine broad domain categories (explicitly listed in Appendix A.5.1):

Basic Operations: FFT, convolution, stencil computation, parallel reduction, matrix multiplication
Graph Algorithms: graph coloring, community detection, strongly connected components, breadth-first traversal
Scientific Computing: fluid simulation, Monte Carlo simulation, sparse linear system solvers, molecular dynamics, heat equation solvers
Machine Learning: k-means clustering, matrix factorization, attention mechanisms, backpropagation, neural network training
Sparse Operations: sparse FFT, sparse tensor operations, sparse convolution, sparse matrix-matrix multiplication
Simulation: cloth simulation, raytracing, N-body simulation, fluid-structure interaction
Image and Signal Processing: feature extraction, image segmentation, video processing, signal transforms, image filtering
Optimization Algorithms: simulated annealing, genetic algorithms, gradient descent, particle swarm optimization
Cryptography and Security: homomorphic encryption, secure hashing, encryption/decryption, blockchain mining, password cracking

Each template contains placeholders like {size}, {optimization}, {algorithm}, {method}, {radius}, etc. For example:

"Generate a CUDA kernel for cloth simulation with a {size}X{size} grid. Optimize for {optimization}."

Variable instantiation. The authors prepare predefined lists of values for each placeholder (full details in Appendix A.5, Table 5). The {size} placeholder is instantiated with values like 32, 64, 128, 256, and 512. The {optimization} placeholder draws from options like "memory bandwidth," "register usage," "multi-GPU scaling," "shared memory," and "warp efficiency." By systematically sampling from these value lists and populating the templates, the pipeline generates a broad range of specific prompts, each requesting a different CUDA kernel with different parameters and optimization targets.

Model choice. The authors use Qwen2.5-Coder 32B (reference [32]) as the generation model, hosted locally and queried via a chat-based API. The choice of a coding-specialized model over a general purpose LLM is driven by the need for compilable, syntactically valid CUDA code — general models are more likely to produce CUDA with subtle syntax errors, missing headers, or incorrect API usage. A 32B parameter coding model strikes a balance between generation quality and computational cost (generating 85k samples requires substantial inference throughput).

Scale and yield. The pipeline generates 85k CUDA samples in total. Of these, only 49.1% compiled successfully, yielding a final set of 46.3k valid files. The 50.9% failure rate comes from three categories of problems the paper identifies: syntactic errors (the model produces code that looks plausible but doesn't parse), missing definitions (the model references functions or types that aren't defined in the generated code), and invalid memory operations (the model writes code that accesses out-of-bounds memory or uses incorrect pointer arithmetic). The high failure rate is not surprising — it reflects the difficulty of generating compilable CUDA code without compiler feedback, a well-known challenge in code generation.

Robustness mechanisms. The pipeline includes fault-tolerance features: retry logic for failed API calls, output validation to check that the model's response actually contains CUDA code (not just explanatory text), file existence checks to handle concurrent access by multiple worker processes, and automatic checkpoint-based resumption so that interrupted generation runs can continue from where they left off.

Why synthetic data matters. The ablation results in Table 4 quantify the contribution: using only Stack (scraped) data yields 17.5% assembly accuracy, while adding synthetic data improves it by +12.5 percentage points (to 30%). The synthetic data provides domain coverage that the scraped data lacks, and more importantly, it introduces patterns of low-level instructions in domains (physics simulation, raytracing, molecular dynamics) that aren't well-represented in public CUDA repositories but are important for real-world GPU workloads.

HIPIFY Transpilation and Its Failure Mode

With CUDA source files in hand (24k scraped + 46.3k synthetic = ~70.3k, plus 6k OpenCL from a separate pipeline), the next step is converting each CUDA file to its HIP equivalent. The authors use AMD's HIPIFY tool (reference [9]), which performs static, rule-based translation of CUDA source to HIP.

What HIPIFY does. HIPIFY works at the source-code level, applying a large set of syntactic substitution rules:

cudaMalloc → hipMalloc
cudaMemcpy → hipMemcpy
cudaFree → hipFree
__global__ → __global__ (unchanged — HIP preserves the kernel launch syntax)
<<<grid, block>>> → hipLaunchKernelGGL(kernel, grid, block, ...)
CUDA-specific types like dim3, cudaStream_t, cudaEvent_t → HIP equivalents

These substitutions are deterministic — HIPIFY is not a learning system, it is a pattern-matching engine with a lookup table of equivalences. This means it handles straightforward API translations reliably but fails on unsupported CUDA features.

The 43.9% failure rate. The paper reports that "files that failed conversion (approx. 43.9%) were discarded" (Section 3.3). This is a substantial loss — nearly half the CUDA files cannot be converted to HIP through HIPIFY alone. The failure modes fall into several categories:

Unsupported CUDA features: CUDA APIs that have no direct HIP equivalent (certain texture memory operations, cooperative groups features, specific versions of dynamic parallelism).
Template instantiation edge cases: CUDA code using complex C++ template metaprogramming where HIPIFY cannot correctly resolve which template expansions correspond to which HIP API calls.
Complex dependency chains: Files that include custom headers defining macros or wrapper functions around CUDA APIs — HIPIFY's pattern matching operates on the preprocessed source and can miss API calls hidden behind custom abstractions.
Version incompatibility: CUDA features introduced in recent CUDA versions that HIPIFY has not yet been updated to support.

What this means for the dataset. The 43.9% failure rate is a filter — the samples that survive have "HIPIFY-compatible" CUDA usage patterns. This introduces a selection bias into the dataset: it overrepresents CUDA code that uses straightforward, well-supported API calls and underrepresents code using advanced or recent CUDA features. The CASS models will therefore learn to translate the kinds of CUDA code that HIPIFY can handle, not arbitrary CUDA. The paper does not discuss this bias explicitly, but it is a direct consequence of using HIPIFY as the alignment mechanism.

Why use HIPIFY at all? The alternative would be manual HIP translation for each sample, which is infeasible at scale (70k samples). Or one could train a model to translate CUDA to HIP directly, but that requires paired training data — the chicken-and-egg problem that CASS is designed to solve. HIPIFY provides a practical, if imperfect, bootstrap: it produces sufficiently many correct translations to build a training set, and the models trained on that set can then outperform HIPIFY itself (the 7B model beats HIPIFY by 7.5% on CASS-Bench, as reported in Section 6).

The Nvidia Compilation Pipeline (Opaque Stack)

Once CUDA-HIP source pairs are available, both must be compiled to extract their assembly representations. The Nvidia pipeline presents a specific challenge: Nvidia's compilation toolchain is opaque — it does not provide direct access to the intermediate steps between source and final binary. The paper characterizes this as a key architectural difference between the two stacks (Figure 2, Section 3.3).

The compilation path. The Nvidia pipeline proceeds through these stages:

nvcc compiles CUDA source (.cu files) into an intermediate representation called PTX (Parallel Thread Execution), which is a virtual ISA that is architecture-independent at the level of NVIDIA GPU generations. PTX is not the final hardware instruction set — it is a portable assembly-like language that gets JIT-compiled to the actual hardware ISA at runtime.
The PTX is then compiled (either at build time or runtime) to SASS (Shader Assembly), the actual binary instruction set for a specific NVIDIA GPU architecture (in this paper: sm_85, corresponding to the Ampere architecture used by the A100).
The SASS binary is embedded into the final executable alongside the host (CPU) code, which is compiled by the host compiler (gcc or similar) into x86-64 assembly.

Why SASS access is difficult. NVIDIA does not provide an official API or command-line flag to extract SASS during compilation. The paper states: "Nvidia provides no access to its binary injection process, device and host assemblies remain intertwined, with no official method for extraction or reintegration" (Section 3.3). To access SASS, the authors must first compile the CUDA code into a complete binary executable, then use cuobjdump (NVIDIA's binary utility, reference [28]) to disassemble the embedded device code back into human-readable SASS. This is a roundabout process: source → compile to binary → disassemble → SASS text.

Host/device separation via regex. Because the Nvidia toolchain interleaves host and device assembly in the final binary, the authors cannot simply extract the device code as a separate file. They develop a regex-based filtering pipeline that parses the combined assembly output and separates host (x86-64) from device (SASS) sections. The specific patterns used are not detailed in the paper beyond "regex-based filtering," but the principle is recognizable to anyone who has worked with GPU binaries: SASS instructions follow specific syntactic patterns (e.g., opcodes like LDG.E, STG.E, FMUL, FADD) that are distinct from x86-64 instructions (movq, call, pushq, jmp), and the boundaries between host and device sections can be identified by looking for cuobjdump-emitted section headers.

Why this matters for the dataset. The ability to separate host and device assembly is critical because CASS's goal is to enable both host-to-host and device-to-device translation independently. A CUDA kernel's host code (which sets up memory, launches the kernel, and copies results) and its device code (which does the parallel computation) are translated through different mechanisms — host code uses x86-64, which is the same on both platforms, while device code uses SASS or RDNA3. Training a single model on the combined assembly would confuse these distinct translation tasks.

Compilation configuration. The authors use the -Os compilation flag (optimize for size) rather than -O3 (optimize for speed). The paper reports that -Os "achieving a 9.3% average token reduction compared to O3" (Section 3.3). This is a practical choice driven by the 16K-token context window constraint: shorter assembly files are more likely to fit within the model's context during training. The size optimization flag produces more compact code (fewer instruction repetitions, less loop unrolling) at the cost of some runtime performance, but for the purpose of training a translation model, the token budget matters more than the benchmark performance of the generated assembly.

Hardware and ISA target. All Nvidia compilation is performed on an A100 PCIe GPU (Ampere architecture). The SASS produced targets the sm_85 ISA, which is specific to the A100's compute capability 8.0. The paper notes (Section 3.3) that while the A100 is a data center GPU, "all CUDA code was compiled targeting the compute capabilities of a standard consumer-grade GPU (e.g., RTX 4090) to maintain parity with the AMD hardware" (Appendix A.2). This means the compilation uses flags that limit the generated SASS to features available on both data center and consumer Ampere GPUs — avoiding instructions that are A100-specific (like certain tensor core operations with particular data layouts).

The AMD Compilation Pipeline (Transparent Stack, Modified)

The AMD compilation pipeline is architecturally different from Nvidia's in ways that are both an advantage and a challenge. The advantage is transparency: AMD's ROCm stack is built on open-source components (LLVM, Clang) that expose the compilation process at every stage. The challenge is that the pipeline must be modified to achieve the host/device separation that CASS requires.

The compilation path. AMD HIP code compiles through these stages (Figure 2, right side):

hipcc (the HIP compiler driver) invokes Clang to compile HIP source into LLVM bitcode (.bc). This is the intermediate representation level, analogous to NVIDIA's PTX.
The LLVM bitcode is then lowered to RDNA3 assembly (the actual hardware ISA for the Radeon RX 7900 XT GPU used in this work) by AMD's LLVM backend.
The critical default behavior (and the modification): In the standard AMD compilation pipeline, the device binary is embedded into the host binary during the transition from bitcode to assembly. The host and device code are combined into a single object file, which is then linked into the final executable. This is convenient for deployment (one file to distribute) but problematic for CASS, which needs independent host and device assembly for translation training.

The modification: deferred device binary insertion. The authors modify this behavior so that device binary insertion is deferred until after host assembly has been converted to object code. The paper describes this as enabling "(1) independent extraction of pure host and device assemblies, and (2) selective recombination for controlled translation and evaluation" (Section 3.3).

In operational terms, the modification likely involves:

Compiling the host code separately from the device code (using Clang's ability to process host and device code independently).
Emitting the device assembly as a standalone file rather than embedding it.
Extracting the host assembly from the host compilation.
Only combining them at the final linking stage (for verification that the translated code is functional).

This is non-trivial engineering — the AMD toolchain is not designed to be used this way, and the authors needed to understand the compilation pipeline well enough to intervene at the right point. The paper does not provide implementation details beyond the conceptual description in Figure 2, but the result is critical: the AMD pipeline can now produce separate host (x86-64) and device (RDNA3) assembly files, exactly matching the structure of the Nvidia pipeline's output after regex-based separation.

Compilation configuration. As with the Nvidia pipeline, the authors use -Os for the AMD compilation, achieving the same ~9.3% token reduction relative to -O3. The AMD compilation targets the RDNA3 ISA, which is the instruction set architecture for AMD's RDNA 3 graphics architecture (used by the RX 7900 XT GPU).

Hardware. All AMD compilation and execution is performed on a system with an Intel i7-14700KF CPU and an AMD Radeon RX 7900 XT GPU (Appendix A.2). The RX 7900 XT is a consumer-class GPU using the RDNA3 architecture, released in late 2022. The paper compiles HIP code targeting this specific GPU's ISA, meaning the generated RDNA3 assembly includes instructions that are specific to this architecture (like s_mov_b32, v_add_co_u32, s_waitcnt — scalar and vector operations unique to AMD's GPU ISA).

Post-compilation filtering. After both stacks have been compiled (CUDA → host x86 + device SASS; HIP → host x86 + device RDNA3), the authors retain only samples that compiled successfully on both pipelines. Asymmetric failures (where one stack succeeds and the other fails) are discarded. The final yield from these steps, including the OpenCL pipeline (discussed next), is ~64k paired samples that form the core of the CASS dataset, before combination with the ~6k OpenCL samples brings the total to 70k.

The OpenCL Pipeline (An Independent Translation Path)

In addition to the CUDA → HIPIFY → HIP path, the authors construct a parallel pipeline that uses OpenCL as the source language (Section 3.4). This is an important methodological choice: OpenCL provides an independent verification that the aligned assembly dataset is not an artifact of the CUDA/HIPIFY conversion process.

What OpenCL provides. OpenCL (Open Computing Language) is a cross-platform, vendor-neutral GPU programming framework maintained by the Khronos Group. Unlike CUDA (Nvidia-only) and HIP (AMD-oriented), OpenCL code can compile and run on both Nvidia and AMD GPUs without source modification. This makes it an ideal "ground truth" source for generating aligned assembly pairs: the same OpenCL kernel, when compiled through each vendor's toolchain, produces functionally equivalent Nvidia and AMD assembly.

The compilation path for OpenCL. The pipeline works differently for each vendor:

On the Nvidia stack: The authors collect approximately 6k OpenCL code snippets from The Stack dataset. Each snippet is compiled using a wrapper C++ function that calls clBuildProgram (the OpenCL API for compiling kernel source, reference [30]), which produces PTX (Nvidia's intermediate representation). The PTX is then processed through the standard Nvidia stack (the same cuobjdump pipeline) to extract SASS device assembly.
On the AMD stack: The same OpenCL files are compiled using Clang directly, which transpiles OpenCL to RDNA3 device assembly while "forcing it to emit intermediate LLVM during this process" (Section 3.4). This path bypasses the HIP toolchain entirely — it goes OpenCL → LLVM IR → RDNA3 assembly, without involving hipcc, HIPIFY, or any HIP-specific components.

Why the OpenCL pipeline matters. The OpenCL samples provide a different distribution of assembly patterns than the CUDA/HIP samples because the translation path is different. OpenCL's memory model, synchronization primitives, and API structure differ from CUDA/HIP, so the generated assembly reflects different programming patterns. The ablation in Table 4 shows that adding OpenCL data improves assembly accuracy by +2.5 percentage points (from 30% to 32.5%), indicating that the OpenCL-sourced assembly provides complementary coverage of the instruction space — patterns that appear in OpenCL-compiled code but not in HIPIFY-converted code.

The OpenCL pipeline also serves as a robustness check: if the CASS models could only translate CUDA/HIP-derived assembly (i.e., if they were learning artifacts of the HIPIFY conversion rather than genuine cross-ISA mapping), they would fail on OpenCL-derived test cases. The fact that OpenCL data contributes positively suggests the models are learning a more general translation capability.

Scale. The OpenCL pipeline contributes ~6k samples to the final dataset, bringing the total to 70,694 samples (Table 2). The combined dataset spans three data sources: Stack-scraped CUDA (24k), synthetic CUDA (46.3k), and OpenCL (6k), with the synthetic data providing the largest fraction by volume and the Stack data providing the most realistic distribution of real-world code patterns.

CASS-Bench Curation (Manually Verified Evaluation)

The CASS-Instruct training dataset is generated automatically through compilation pipelines, which means it contains noise — samples that compile and produce assembly but may have subtle semantic errors (off-by-one arithmetic, incorrect boundary conditions, uninitialized memory) that the compilation process doesn't detect. To provide a clean evaluation signal, the authors curate CASS-Bench, a manually verified benchmark of 40 samples across 16 GPU-centric domains (Section 4.2).

The curation protocol. The process is iterative and involves human verification at multiple stages:

Generate CUDA implementations. For each of the 40 prompts (each prompt is a natural language description of a GPU computation task), the authors use Claude-3.7 to generate a CUDA implementation. Claude-3.7 is chosen as the generation model presumably because it produces high-quality, compilable CUDA code with fewer errors than smaller models — the goal is to minimize the number of iterations needed to reach a correct implementation.
Compile and execute on Nvidia hardware. Each generated CUDA implementation is compiled and run on the Nvidia A100 GPU. The output of the computation (e.g., a numerical result, a transformed array, a simulation state) is captured as the reference ground truth.
Generate corresponding AMD code. The authors prompt Claude-3.7 to generate the AMD-equivalent code — either HIP source (for source translation tests) or RDNA3 assembly (for assembly translation tests).
Verify output equivalence. The generated AMD code is compiled and executed on the AMD RX 7900 XT GPU. Its output is compared against the Nvidia reference output. If the outputs mismatch, the AMD code is regenerated.
Iterative correction. Mismatches can arise from three sources: compilation errors (the generated AMD code doesn't compile), formatting differences (the output is numerically correct but formatted differently — different precision, different whitespace, different ordering), or random generator variance (the code uses random number generators that produce different sequences on Nvidia vs. AMD hardware). For formatting and random variance issues, the authors adjust the AMD code to use the same random seed and same formatting conventions as the CUDA version. For compilation errors, they regenerate the AMD code with additional prompting to fix the specific error.
Manual verification. Only samples where a human verifier confirms "output equivalence" after manual inspection are included in CASS-Bench. This means CASS-Bench is not just measuring whether the translated code compiles — it measures whether it produces exactly the same computational result as the original CUDA code when executed on AMD hardware.
Assembly extraction. Once a CUDA-HIP pair passes the equivalence check, both the original CUDA source and the verified HIP source are run through the CASS compilation pipeline (Section 3.3) to extract aligned host and device assembly — producing the same format as the training data but with execution-verified correctness.

Domain coverage. The 40 samples span 16 domains, as shown in Figure 3 (right). Each domain is represented by 1–5 curated prompts. The domains include categories that appear in the training data distribution (linear algebra, machine learning, physics simulation) as well as categories that are intentionally distinct to test generalization (graph algorithms, data structures, cryptography). The domain distribution is designed to be broad enough to reveal which types of GPU computation the model handles well and which it struggles with.

Why manual verification matters. The training dataset is large but noisy — some fraction of the 70k samples likely contain subtle bugs that survived compilation. CASS-Bench provides a clean evaluation signal: when the model produces a translation that passes the benchmark, it means the translation is not just syntactically plausible but functionally correct. The paper's finding that assembly accuracy drops to 0% on math, data structures, and graph tasks in CASS-Bench (Section 6, Figure 5) is known to be a real capability gap, not an artifact of noisy evaluation data.

Model Training: Supervised Fine-Tuning of Qwen2.5-Coder

With the CASS-Instruct dataset in hand, the authors fine-tune language models for the translation task. The base model is Qwen2.5-Coder (reference [32]), pre-trained on a large corpus of code (including, importantly, some CUDA and HIP code — though without the paired alignment that CASS provides). The choice of a code-specialized base model over a general language model matters because the base model already knows the syntax of CUDA and HIP — it recognizes cudaMalloc as a function call, __global__ as a kernel qualifier, and <<<>>> as kernel launch syntax. Fine-tuning on CASS teaches it the mapping between these syntactic elements, not the syntax itself.

Model scales. Three variants are trained: 1.5B, 3B, and 7B parameters. Training all three scales serves two purposes: it enables ablation on model size (how much does scale matter for assembly translation?), and it provides smaller, more deployable models for practitioners who cannot run a 7B model.

Two task variants. The authors train separate models for source translation and assembly translation. The source translation model takes CUDA source as input and outputs HIP source. The assembly translation model takes SASS assembly (either host x86 or device SASS) as input and outputs RDNA3 assembly (or host x86, for the host translation task). The paper does not specify whether a single model handles both host and device translation or whether separate models are trained for host and device, but the phrasing in Section 5 suggests a single assembly model that handles both: "Two variants are developed: one for assembly translation and another for source translation."

Input normalization (assembly-specific). Assembly code presents a token efficiency challenge: it contains significant whitespace (indentation, alignment of operands) and varying amounts of comments (disassembly tools often annotate each instruction with address offsets, hex encodings, and human-readable comments). Since the training context window is limited to 16K tokens, this padding wastes capacity. The authors apply different normalization strategies to the two architectures:

CUDA (SASS) assembly: "normalized... by removing redundant whitespace and comments, which reduced token count by roughly 15%" (Section 5). This is a careful choice — SASS syntax is relatively robust to whitespace changes (instructions are delimited by newlines and semicolons, not by indentation), so removing whitespace doesn't change the parse.
HIP (RDNA3) assembly: "No preprocessing was applied to HIP assembly code due to its sensitivity to whitespace changes" (Section 5). Some RDNA3 assembly constructs may use whitespace as a significant delimiter (e.g., certain macro expansions or multi-line directives), so aggressive whitespace stripping could produce invalid assembly.

Training hyperparameters. The full configuration (Section 5):

Hardware: 4× A100 GPUs (the same A100 used for compilation, on the Nvidia server).
Batch size: 4 per GPU.
Gradient accumulation: 32 steps.
Effective batch size: 4 × 4 × 32 = 512. This is a large effective batch size, which helps stabilize training when the dataset has diverse, sometimes noisy samples.
Learning rate: $1 \times 10^{-5}$ . The paper describes this as "relatively aggressive" and justifies it by noting the dataset's "distributional divergence from the models' pretraining corpus" (Section 5). In plain language: the CASS dataset looks very different from the code on which Qwen2.5-Coder was pretrained (which is mostly high-level source code, not assembly), so the model needs larger weight updates to adapt to the new distribution. A smaller learning rate (e.g., $5 \times 10^{-6}$ , typical for instruction tuning) would converge too slowly given this distribution shift.
Context window: 16K tokens during training.
Inference context extension: At inference time, the authors apply RoPE extrapolation to support up to 32.7K tokens. RoPE (Rotary Position Embedding, reference [37]) is the positional encoding scheme used by Qwen2.5-Coder. Extrapolation means using the model with sequences longer than it was trained on by scaling the position indices — the sine/cosine frequencies that encode position information are stretched to cover the longer range. This is not guaranteed to work well (attention patterns may degrade at unseen positions), but the paper's results suggest it works well enough for the assembly domain, where longer sequences tend to be repetitive (more instructions, but the same kinds of instructions in the same patterns).

Training infrastructure. The paper uses three optimization components:

DeepSpeed (reference [33]) with optimizer state sharding (ZeRO Stage 2). This distributes the optimizer states (AdamW moment estimates) across GPUs, reducing per-GPU memory consumption and enabling larger batch sizes.
Liger Kernel (reference [34]), which provides optimized Triton kernels for common LLM training operations (cross-entropy loss, fused linear layers), reducing training time.
Paged AdamW optimizer (reference [35]), which uses a paged memory management scheme for optimizer states, further reducing memory pressure.

The combination achieves "98% GPU utilization" (Section 5), indicating efficient use of the available hardware — the GPUs are computing almost continuously rather than waiting for data loading or communication.

Training framework. All fine-tuning is implemented using LLaMA-Factory (reference [36]), a unified framework that wraps model loading, data processing, training loop, and evaluation. Using LLaMA-Factory rather than custom training code reduces engineering overhead and makes the training recipe reproducible.

Inference cost. The paper reports that "inference was efficient, requiring approximately 56 seconds per a 16K-token sample" (Section 5). This is on the 4× A100 setup, presumably with the 7B model. At this rate, evaluating all 40 CASS-Bench samples takes roughly 37 minutes — practical for research but non-trivial for production deployment.

Training duration. Figure 10 (Appendix) shows accuracy vs. training steps for all three model scales. The curves show that assembly accuracy improves gradually over training steps, with the 7B model reaching its peak around 600–800 steps, while the smaller models saturate earlier (around 400–600 steps). The paper does not specify the total number of training steps or epochs, but the effective batch size of 512 applied to 70k samples means approximately 137 steps per epoch (70,694 / 512 ≈ 138), suggesting the models are trained for roughly 4–6 epochs.

Summary of Design Choices and Their Justifications

Repository-level scraping over file-level scraping: preserves dependencies and build configurations, improving compilation success by 23.7% — without this, the scraped data pipeline would yield far fewer compilable samples.
Synthetic data generation with Qwen2.5-Coder 32B over smaller models or no synthetic data: fills domain coverage gaps in the scraped data and contributes +12.5% assembly accuracy in ablation — the persona-augmented prompting strategy ensures diversity across 9 domain categories.
HIPIFY for CUDA-to-HIP alignment over manual translation: the only scalable option for 70k samples, despite the 43.9% failure rate — accepts selection bias in exchange for volume, and the trained models can then outperform HIPIFY itself.
-Os compilation over -O3: prioritizes token efficiency (9.3% reduction) to fit within 16K context windows over runtime performance optimization — a practical constraint, not a quality claim.
Modified AMD compilation pipeline (deferred device binary insertion): enables extraction of independent host and device assembly, which the default AMD toolchain does not support — this is the key engineering contribution that makes assembly-level dataset construction possible.
OpenCL pipeline as an independent data source: provides complementary assembly patterns that don't go through HIPIFY, contributing +2.5% accuracy and serving as evidence that the model's translation capability is not an artifact of the CUDA/HIPIFY path.
CASS-Bench manual verification: provides a clean evaluation signal with execution-proven ground truth, avoiding the noise inherent in the automatically constructed training set — essential for measuring progress on functional correctness, not just syntactic plausibility.
Aggressive learning rate ( $1 \times 10^{-5}$ ): justified by the distribution shift between the CASS dataset (assembly-heavy) and the base model's pretraining corpus (source-code-heavy) — the model needs larger weight updates to adapt to the new domain.
RoPE extrapolation for inference: extends the trained 16K context window to 32.7K tokens at inference time, enabling the model to handle longer assembly sequences than it saw during training — practical for real assembly files that may exceed the training context length.
Assembly-specific preprocessing (whitespace removal for CUDA, no preprocessing for HIP): reflects architectural differences in ISA syntax — SASS is robust to whitespace changes, while RDNA3 may use whitespace-significant constructs.
Three model scales (1.5B, 3B, 7B): enables analysis of scaling behavior and provides deployment flexibility — even the 1.5B model (90% source accuracy, 17.5% assembly accuracy) substantially outperforms all commercial baselines on both tasks.

4. Key Insights and Innovations

Innovation 1: The Dataset IS the Contribution — Reframing GPU Transpilation as a Data Availability Problem Rather Than a Translation Algorithm Problem

The paper's most fundamental intellectual move is its diagnosis of why GPU-to-GPU assembly translation hasn't been tackled before. The standard assumption in software portability is that the bottleneck is algorithmic — that we need a smarter translator, a better intermediate representation, or a more sophisticated static analysis. CASS argues that this diagnosis is wrong, or at least incomplete: the real bottleneck is the complete absence of aligned, paired training data across GPU instruction set architectures. Without such data, no translation algorithm — whether rule-based, IR-level, or machine-learned — can be developed or evaluated.

This reframing is significant because it shifts the problem from one that looks like a compiler engineering challenge (design a better transpiler) to one that looks like a data infrastructure challenge (build a pipeline that can generate aligned pairs at scale). The prior work that the paper cites — HIPIFY, CuPBoP-AMD, ZLUDA, GPU Ocelot — all attempted the former. They built increasingly sophisticated translation mechanisms operating at different levels of the compilation stack (source, IR, binary), but all hit fundamental ceilings: HIPIFY fails on ~44% of CUDA files (Section 3.3); ZLUDA achieves only 2.5% assembly accuracy on CASS-Bench (Table 3) despite operating directly on compiled binaries; GPU Ocelot was abandoned due to scalability and overhead issues. The paper's implicit argument is that these failures are not primarily algorithmic — they stem from each tool's developers having to manually encode translation rules without access to a systematic corpus of aligned examples that would reveal the full scope of the mapping problem.

What makes this reframing more than a semantic distinction is that it redefines the measure of progress. In the compiler-engineering frame, progress means a higher success rate on an ad-hoc test suite. In the data-infrastructure frame, progress means a larger, more diverse, better-aligned dataset — and the translation accuracy follows from the data quality. The ablation results in Table 4 make this argument empirically: Stack data alone gives 17.5% assembly accuracy; adding synthetic data adds +12.5%; adding OpenCL adds another +2.5%; RoPE extrapolation adds +5%. Each data contribution directly improves the model, with no change to the translation algorithm. This suggests that the marginal return to better data is high, and that the field's historical focus on better algorithms may have been premature when the training corpus was essentially empty.

The comparison to CPU assembly translation is instructive here. CRT (Heakl et al., 2024) and Guess & Sketch (Lee et al., 2023) demonstrated that language models can translate between CPU ISAs given aligned training data. The field might have concluded that the same approach would work for GPUs — but nobody tried, because the data didn't exist. CASS's insight is that the CPU precedent was actually a proof of concept for a data-first approach, not an algorithmic innovation waiting to be ported. The paper's contribution is to recognize and fill the data gap that was blocking the port.

There is a subtlety worth surfacing: this reframing is partly a rhetorical move that justifies the paper's resource investment. Building a 70k-sample aligned dataset across two GPU stacks required substantial engineering (repository scraping, synthetic generation, dual-stack compilation, host/device separation) but relatively modest intellectual novelty in the individual components. By arguing that the dataset is the contribution, the paper positions this engineering effort as the primary intellectual contribution rather than as infrastructure in service of a model innovation. This is defensible — the dataset enables future work that was previously impossible — but it also means the paper's significance depends heavily on whether the community adopts CASS as infrastructure. If it does, the reframing becomes self-fulfilling: CASS will have made GPU assembly translation a data-driven field. If it doesn't, the paper remains a one-off demonstration that domain-specific fine-tuning helps on a niche task.

Innovation 2: The Dual-Stack Compilation Pipeline as a General Method for Generating Aligned Low-Level Code Pairs

Beyond the specific dataset, the paper contributes a methodological template for generating aligned assembly pairs across any pair of GPU (or, in principle, any accelerator) architectures that share a common high-level programming model. This is not an incremental improvement on existing dataset construction techniques — it is a novel strategy that the paper invents to solve a problem that had no prior solution.

The key insight is that semantic alignment can be guaranteed by construction if both assembly files are compiled from the same algorithm specification, rather than relying on post-hoc matching or manual annotation. The pipeline operates as follows: start with a single CUDA source file → convert it to HIP via HIPIFY → compile both through their respective vendor toolchains → extract the resulting assembly. Because both assembly files originate from the same computational specification (the algorithm expressed in CUDA, mechanically translated to HIP), they are guaranteed to be semantically equivalent — they compute the same function, even if the instruction sequences look very different.

This is conceptually different from how paired datasets are typically constructed in machine translation, where human translators produce parallel corpora, or in code translation, where aligned pairs are scraped from repositories that contain both versions (e.g., a project that maintains both a Python and a C++ implementation). Those approaches require the alignment to exist before dataset construction. The CASS approach creates the alignment through compilation — it takes a single codebase and produces paired assembly by routing it through two different toolchains. This is possible because GPU programming models (CUDA, HIP, OpenCL) are designed to abstract over hardware differences, providing a "pivot language" that both vendors' compilers can lower to their respective ISAs.

The methodological contribution has implications beyond Nvidia-to-AMD translation. The same strategy could generate aligned pairs for:

Different GPU generations from the same vendor (e.g., Nvidia SASS sm_80 vs. sm_90), enabling translation between GPU architectures that share a vendor but have different ISAs.
Intel GPUs (which use the oneAPI/Level Zero programming model), extending the cross-vendor coverage beyond Nvidia and AMD.
Specialized accelerators (TPUs, NPUs, FPGAs) as long as there exists a high-level programming model (like OpenCL or SYCL) that compiles to both the accelerator and a reference GPU architecture.
CPU ISAs using a portable language like C as the pivot — compile the same C code with gcc -S for x86 and clang -S for ARM to get aligned pairs without manual translation.

The paper does not make these generalizations explicitly, but the pipeline design is clearly intended as a template rather than a one-off artifact. The OpenCL pipeline (Section 3.4) is the first demonstration of this generality: by using OpenCL as the pivot language instead of CUDA/HIP, the same alignment strategy works without going through HIPIFY at all. The fact that OpenCL-sourced data adds +2.5% to assembly accuracy (Table 4) confirms that the approach generalizes beyond the specific CUDA→HIPIFY→HIP path.

The most technically non-trivial aspect of this contribution is the host/device separation engineering. As detailed in Section 3.3 and Figure 2, the Nvidia and AMD compilation stacks have fundamentally different architectures regarding how host and device code are combined. Nvidia's stack is opaque — host and device assembly are intertwined in the final binary, and the only way to access device SASS is to compile to binary and then disassemble. AMD's stack is transparent but embeds the device binary into the host during bitcode-to-assembly transition, requiring pipeline modification to extract them separately. The paper solves both problems — regex-based separation for Nvidia, deferred device insertion for AMD — creating a unified output format (separate host and device assembly files for each vendor) from two fundamentally different compilation architectures. This engineering is what makes the methodological template practical: future efforts targeting different GPU pairs will face similar host/device separation challenges, and the paper's solutions provide a starting point.

Innovation 3: Assembly Translation Accuracy as a Domain-Dependent, Not Uniform, Property — With Zero-Shot Failure as the Diagnostic Signal

The paper's most striking empirical finding is not the 37.5% assembly accuracy per se, but the extreme heterogeneity of that accuracy across domains — from 0% in math, data structures, and graph tasks to 25–50% in linear algebra and memory operations, up to 100% in physics simulations (Section 6, Figure 5, and Appendix A.4.2). This heterogeneity is not a weakness of the approach; it is a diagnostic discovery about the nature of GPU assembly translation that was not obvious before the CASS dataset and benchmark existed.

The finding reveals that learnability of assembly translation is domain-dependent in a way that source translation is not. At the source level, the CASS models achieve 90–95% accuracy uniformly across domains (Figure 5, left panel) — the syntactic mapping from CUDA to HIP API calls is largely domain-independent because it is a surface-level syntactic transformation. At the assembly level, the mapping from SASS to RDNA3 depends on the semantic patterns of the computation: physics simulations (which tend to use simple, repetitive control flow and regular memory access patterns) produce assembly that the model can learn to translate, while graph algorithms and data structure operations (which involve irregular memory access, pointer chasing, and complex control flow) produce assembly patterns that the model completely fails to translate.

This is a conceptual finding about the structure of GPU ISAs, not merely a performance report. It tells us that the mapping between SASS and RDNA3 is not a uniform function — it is a function whose complexity varies with the computational pattern being expressed. For regular, loop-heavy, floating-point-intensive computations (the "easy" domains), the two ISAs express the same computation in systematically related ways — perhaps a loop over array elements in SASS becomes a predictable sequence of vector loads and arithmetic in RDNA3. For irregular, control-flow-heavy computations, the mapping is apparently too complex for the current model to learn from 70k examples — the relationship between SASS and RDNA3 instructions is too context-dependent, too sensitive to register allocation decisions and instruction scheduling, for a sequence-to-sequence model to capture without more data or architectural innovation.

The 0% accuracy cases are particularly informative because they function as a diagnostic of insufficient representation — either insufficient training data in those domains or insufficient model capacity to learn the mapping. The paper's domain distribution (Figure 3, left) shows that math, data structures, and graph algorithms are modestly represented in the training set, but apparently not enough to achieve non-zero translation accuracy. This suggests that assembly translation may require orders of magnitude more data in irregular-computation domains than in regular-computation domains — a scaling property that was not predictable a priori and that has implications for how future dataset collection efforts should prioritize domain coverage.

This finding also refines the paper's headline claim in an intellectually honest way. The 37.5% assembly accuracy is an average that hides variance from 0% to 100%. A reader focused only on the average might conclude that assembly translation is "promising but not ready." The domain-level breakdown reveals a more nuanced picture: assembly translation works surprisingly well for some computational patterns and fails completely for others, suggesting that the path to higher overall accuracy is not simply "more data" but "more data in specific underrepresented domains" — a much more actionable and specific research direction.

Innovation 4: The Conceptual Separation of Source and Assembly Translation as Distinct Problems with Distinct Difficulty Profiles

The paper makes a taxonomic contribution that is easy to overlook: it establishes source-level and assembly-level GPU translation as fundamentally different problem classes with different difficulty characteristics, different failure modes, different baseline capabilities of general-purpose models, and different scaling behaviors. This separation was not obvious before CASS because no prior work had attempted both tasks with the same models on the same benchmark under controlled conditions.

The evidence for the distinctness of the two tasks is stark. At the source level, general-purpose models succeed without domain-specific training: GPT-4o achieves 80% accuracy, Claude-3.7 achieves 85%, and Qwen2.5-Coder-32B achieves 87.5% (Table 3, source translation columns). This means that pretraining on a large corpus of source code (which includes both CUDA and HIP files, just not aligned pairs) is sufficient to learn the CUDA→HIP mapping — the mapping is shallow enough that exposure to both languages in unaligned form captures most of the required knowledge.

At the assembly level, every general-purpose model achieves 0% accuracy — GPT-4o, Claude-3.7, Gemini 2.0 Flash, Qwen2.5-Coder-32B, and ZLUDA all score 0% or near-0% (Table 3, assembly translation columns). This means that pretraining on code corpora (which include assembly files — The Stack contains x86 assembly, for example) is not sufficient to learn the SASS→RDNA3 mapping. The mapping is deep enough — it involves instruction-level correspondences, register allocation, memory addressing modes, and synchronization semantics — that it requires explicit paired training data.

This contrast has a clear implication: source translation is a solved problem for general-purpose models at current accuracy targets (80–95%), while assembly translation is impossible for general-purpose models and requires domain-specific fine-tuning. The paper does not state this implication in exactly these terms, but the data in Table 3 supports it. Even the smallest CASS model (1.5B, achieving 90% source accuracy) matches or exceeds all proprietary models on source translation, suggesting that source translation capability is limited mainly by the base model's exposure to CUDA/HIP syntax, not by model scale or specialized training. Assembly translation, by contrast, shows a clear scaling trend with model size (17.5% → 25% → 37.5% for 1.5B → 3B → 7B; Table 3), suggesting that larger models can capture more of the complex SASS→RDNA3 mapping given sufficient paired data.

This separation is not merely descriptive — it has practical implications for how GPU translation systems should be architected. A production system might use a small, cheap model for source translation (where accuracy is already high) and reserve a larger, domain-specialized model for assembly translation (where every parameter counts). Or it might use the source-level model as a first pass and only invoke the assembly-level model for performance-critical kernels where low-level optimizations matter. The paper does not explore these architectural implications, but the clear separation it establishes between the two tasks provides the empirical foundation for such exploration.

Innovation 5: Execution Fidelity as a Co-Emergent Property — The Finding That Translation Accuracy and Runtime/Memory Preservation Are Correlated

The paper reports a finding that seems almost too good to be true, and its significance lies precisely in the fact that it is true (within the measured limits): "over 85% of samples fall within ±5.6% across both metrics" (memory usage and execution time, Section 6 and Figure 11), meaning that when the CASS model produces assembly that compiles and produces correct output, it also tends to produce assembly with runtime and memory characteristics very close to the native, hand-written HIP equivalent.

This is not an obvious consequence of the training objective. The model was trained to maximize token-level accuracy — it learns to predict the next RDNA3 instruction given the preceding SASS instructions. There is nothing in the loss function about runtime, memory usage, register pressure, or instruction-level parallelism. A model optimizing purely for token accuracy could, in principle, produce "correct" assembly that uses more instructions, spills more registers to memory, or issues more redundant loads — functionally correct but performance-degraded. The fact that it largely does not suggests that correctness and efficiency are correlated in the assembly space in a way that sequence-to-sequence learning can exploit: the token-level mapping from SASS to RDNA3, when learned accurately, tends to preserve the computational structure that determines performance.

This is a subtle point that the paper does not fully unpack, but it has implications for how we think about ML-based compilation. If translation accuracy and performance preservation were uncorrelated, we would need separate objectives — one for correctness, another for performance — and likely separate evaluation pipelines (compilation correctness tests plus performance benchmarks). The correlation means that accuracy on CASS-Bench is a reasonable proxy for both correctness and efficiency, which dramatically simplifies evaluation for future work. Researchers can measure translation accuracy and have reasonable confidence that accuracy gains will translate (roughly) to performance preservation, without running full benchmark suites for every model checkpoint.

There is a caveat: the correlation may be partly an artifact of the -Os compilation flag, which the authors chose specifically to reduce code size. -Os tends to produce more uniform, less aggressively optimized assembly than -O3, potentially making the SASS→RDNA3 mapping more regular and the performance characteristics more predictable. Under -O3, where the two compilers might make very different optimization decisions (different loop unrolling factors, different vectorization strategies, different inlining decisions), the correlation between token-level accuracy and runtime performance might weaken. The paper does not explore this, and it represents a limitation of the current evidence — the 85%-within-±5.6% finding applies to -Os-compiled code and may not generalize to performance-optimized compilation.

Nonetheless, the finding matters because it lowers the barrier to entry for future work on GPU assembly translation. A team wanting to extend CASS to new GPU architectures or new domains can evaluate their model primarily on CASS-Bench accuracy and have some confidence that accuracy improvements will translate to usable code, without needing to set up the full hardware testing pipeline (which requires access to both Nvidia and AMD GPUs) for every experiment. This is a pragmatic contribution — it makes the research more accessible — with theoretical undertones about the relationship between syntactic correctness and semantic performance in learned compiler output.

5. Experimental Analysis

Evaluation Methodology

Dataset. The primary training dataset is CASS-Instruct, comprising 70,694 samples spanning scraped CUDA from The Stack v2 (24k), synthetically generated CUDA (46.3k), and OpenCL (6k), all compiled through the dual-stack pipeline to produce aligned CUDA/HIP source and SASS/RDNA3 assembly pairs. The evaluation dataset is CASS-Bench, a manually curated suite of 40 samples across 16 GPU-centric domains, each execution-verified to ensure that the Nvidia and AMD outputs are functionally equivalent. All experiments use the same CASS-Bench split for both source and assembly evaluation, with no separate train/test split needed since CASS-Bench was constructed independently of CASS-Instruct.
Base model(s). The authors fine-tune Qwen2.5-Coder at three parameter scales: 1.5B, 3B, and 7B parameters. Qwen2.5-Coder is a code-specialized language model pretrained on a large corpus of source code, including CUDA and HIP files in unaligned form. The choice of a code-specialized base model over a general-purpose LLM is motivated by the need for syntactic knowledge of CUDA and HIP before fine-tuning — the base model already recognizes GPU programming constructs, and the CASS dataset teaches the cross-architecture mapping rather than the syntax itself. Two task-specific variants are trained: one for source translation (CUDA → HIP) and one for assembly translation (SASS → RDNA3 for device code, plus corresponding host x86 translation).
Metrics. The primary metric is translation accuracy, defined as the fraction of CASS-Bench samples for which the model-generated output (HIP source or RDNA3 assembly) compiles successfully and produces the correct computational output when executed on AMD hardware, matching the Nvidia reference output. This is a strict functional correctness metric — it requires both compilation success and output equivalence, not merely syntactic plausibility. Secondary metrics include memory usage deviation and execution time deviation relative to native HIP baselines, measured in percentage terms. Each test is executed 20 times with reported values reflecting the average across runs "to mitigate noise and ensure statistical reliability" (Appendix A.4.2).
Baselines. The paper evaluates against five categories of baselines, listed in Table 3:
- Proprietary LLMs: GPT-4o (Hurst et al., 2024, reference [12]), Claude-3.7 Sonnet (Anthropic, February 2025, reference [13]), and Gemini 2.0 Flash (Hassabis and Kavukcuoglu, December 2024, reference [38]). These represent state-of-the-art general-purpose models with no domain-specific fine-tuning on GPU translation.
- Open-source code LLMs: Qwen2.5-Coder-32B (the base model without CASS fine-tuning, reference [32]) and Qwen2.5-Coder-7B-Instruct (an instruction-tuned variant). These test whether general code pretraining alone, without paired GPU data, suffices for assembly translation.
- Static translation tool: HIPIFY (AMD, 2025, reference [9]), the production rule-based CUDA-to-HIP source translator. This is the most directly comparable baseline for source translation, representing the state of the art in non-ML approaches.
- Runtime-level system: ZLUDA (Janik, 2024, reference [6]), which dynamically translates CUDA binaries to AMD-compatible code. ZLUDA is evaluated at the assembly level as the closest existing system to assembly translation, operating by intercepting CUDA runtime calls and translating embedded PTX/SASS into AMD-compatible LLVM IR.
- CASS model family: The authors' fine-tuned models at 1.5B, 3B, and 7B scales, evaluated in both source and assembly translation configurations with and without RoPE extrapolation.
ZLUDA evaluation deserves specific methodological detail (described in Appendix A.1): the authors design a two-track strategy. In the source-to-source setting, they compile CUDA source into PTX using nvcc, then feed PTX to ZLUDA for translation into AMD-compatible LLVM IR and lowering to RDNA3. In the assembly-to-assembly setting, they compile CUDA into a complete executable, invoke it directly, and ZLUDA intercepts CUDA runtime calls, dynamically translating embedded PTX/SASS before execution on the AMD backend. This dual strategy tests both ZLUDA's static and runtime translation capabilities.
Generation budget / compute accounting. All models generate a single output per prompt (greedy decoding or sampling with temperature not explicitly specified, but the paper describes deterministic instruction-following behavior), meaning there is no test-time compute scaling — the comparison is between models at their default generation settings, not between strategies for spending a generation budget. For fair comparison across models, the metric is simply whether the single generated output achieves functional equivalence. The paper measures inference cost at "approximately 56 seconds per a 16K-token sample" for the 7B model on 4× A100 GPUs (Section 5), establishing a practical estimate of deployment cost. The training compute is not explicitly quantified in FLOPs, but the paper reports GPU utilization (98%) and hardware configuration (4× A100) to enable rough reproduction.
Cross-validation / statistical protocol. No formal cross-validation or statistical significance testing is reported. The evaluation is a single-pass measurement of all models on all 40 CASS-Bench samples. For the runtime and memory measurements, each test is executed 20 times and results are averaged "to mitigate noise and ensure statistical reliability" (Appendix A.4.2), but no confidence intervals, standard deviations, or significance tests are reported — only percentage ranges. The ablation study (Table 4) reports point estimates without error bars. The paper does not address whether the 40-sample CASS-Bench test set is large enough for statistically meaningful comparisons — a limitation given that some domain categories contain only 1-2 samples and differences between models (e.g., CASS-7B at 37.5% vs. Qwen2.5-Coder-32B at 25%) represent only a small number of absolute correct predictions.

Main Quantitative Results

Assembly-to-Assembly Translation: CASS Models Achieve Non-Zero Accuracy Where All Baselines Score 0%

Headline result. The 7B CASS model achieves 37.5% assembly translation accuracy on CASS-Bench (Table 3), while GPT-4o, Claude-3.7, and Gemini 2.0 Flash all score 0%, and Qwen2.5-Coder-32B (the base model without CASS fine-tuning) achieves only 25% (but only after the model was "explicitly prompted," a detail that is not elaborated). The full results from Table 3:

Model	Assembly Accuracy (%)
GPT-4o	0
Claude-3.7	0
Gemini 2.0 Flash	0
Qwen2.5-Coder-32B	25.0
Qwen2.5-Coder-7B-Instruct	0
ZLUDA	2.5
CASS-1.5B	17.5
CASS-3B	25.0
CASS-7B	37.5

The 0% accuracy of all general-purpose models is the most informative result. GPT-4o, Claude-3.7, and Gemini 2.0 Flash have all been trained on vast corpora that include assembly code (x86, ARM, and potentially some GPU assembly through documentation and repositories), yet none can translate even a single SASS sample to functionally correct RDNA3 assembly. This confirms the paper's central premise: pretraining on unaligned code corpora does not impart GPU assembly translation capability. The contrast with the 25% achieved by the base Qwen2.5-Coder-32B (without CASS fine-tuning) is notable but requires the caveat that the paper states this result was obtained after "explicitly prompted" — likely requiring carefully engineered prompts that specify the translation task in detail, unlike the simpler prompts used for the CASS models. The Qwen2.5-Coder-7B-Instruct (instruction-tuned variant) achieving 0% suggests that instruction tuning on general code tasks may actually interfere with assembly translation capability, perhaps by training the model to refuse unfamiliar tasks or to produce explanatory text rather than raw assembly.

ZLUDA's 2.5% accuracy. ZLUDA is a qualitatively different baseline because it is not a language model — it is a runtime translation system that intercepts CUDA API calls and dynamically translates PTX/SASS to AMD-compatible code. Its 2.5% accuracy (representing exactly 1 of the 40 CASS-Bench samples) is attributed by the authors to "its compatibility with RDNA1" (Section 6, discussion of Table 3), since ZLUDA was originally developed for RDNA1 and may not fully support RDNA3 instructions. This result is somewhat misleadingly presented: ZLUDA's translation mechanism operates at the LLVM IR level (not the hardware assembly level), so it is not directly comparable to the CASS models' assembly-to-assembly task. ZLUDA's near-zero score is less a reflection of its translation capability and more a reflection of the benchmark task exceeding its design scope.

Scaling behavior with model size. Assembly accuracy scales monotonically with model size: 1.5B → 17.5%, 3B → 25.0%, 7B → 37.5% (Table 3). This is a clean scaling trend suggesting that larger models can capture more of the complex SASS→RDNA3 mapping. The paper does not test models larger than 7B (presumably due to computational constraints of 4× A100 GPUs with 16K-token context windows), so the scaling trend cannot be extrapolated beyond this range.

Domain-dependent accuracy breakdown (Figure 5, right panel). The 37.5% average masks extreme heterogeneity across domains:

0% accuracy: Math, data structures, and graph algorithm tasks.
25–50% accuracy: Linear algebra and memory operations.
100% accuracy: Physics simulation tasks.

The paper does not report per-domain sample counts, but Figure 3 (right) shows the CASS-Bench domain distribution with 1–5 samples per domain. A single domain having 100% accuracy on its 1–5 samples is far less meaningful than a domain with 50% accuracy on 5 samples, but the paper treats all domain-level accuracies as equivalent. The Appendix (A.4.2) provides slightly more detail: "assembly accuracy is inconsistent, 0% in Math, Data Structures, and Graph, 25–50% in linear algebra and memory operations." The 100% for physics simulation is attributed to "simpler or repetitive control flows" — an interpretation consistent with the paper's broader claim that regular, floating-point-heavy computations are easier to translate than irregular, control-flow-heavy computations.

Source-to-Source Translation: CASS Models Outperform All Baselines, Including Proprietary Models 100× Larger

Headline result. The 7B CASS model achieves 95% source translation accuracy on CASS-Bench (Table 3), while the best commercial baseline (Claude-3.7) achieves 85%, the best open-source baseline (Qwen2.5-Coder-32B) achieves 87.5%, and HIPIFY achieves 87.5%. Full results from Table 3:

Model	Source Accuracy (%)
GPT-4o	80.0
Claude-3.7	85.0
Qwen2.5-Coder-32B	87.5
HIPIFY	87.5
CASS-1.5B	90.0
CASS-3B	92.5
CASS-7B	95.0

Comparison to HIPIFY. The paper emphasizes that the CASS-7B model "surpassed HIPIFY by 7.5%" (Section 6), but this requires careful interpretation. HIPIFY's 87.5% accuracy represents its performance on the 40 CASS-Bench samples. These samples are not representative of arbitrary CUDA code — they were generated by Claude-3.7 from natural language prompts and manually verified, which means they are syntactically clean, well-structured, and avoid the kinds of unsupported CUDA features that cause HIPIFY's high failure rate (~43.9% in the training pipeline). In other words, the 87.5% is HIPIFY's accuracy on code that HIPIFY can already handle well — the CASS-Bench samples implicitly filter out the hardest cases that cause HIPIFY to fail entirely. The 7.5% improvement therefore represents the model's ability to handle edge cases within the "HIPIFY-compatible" subset of CUDA, not its ability to translate code that HIPIFY cannot handle at all (since those samples would not appear in CASS-Bench). The paper does not discuss this selection effect.

Performance of general-purpose models. All three proprietary models achieve non-zero source translation accuracy (80–85%), demonstrating that general-purpose code pretraining can learn the CUDA→HIP mapping to a substantial degree without paired training data. This is the key contrast with assembly translation: source translation is learnable from unaligned corpora, while assembly translation is not. GPT-4o's 80% accuracy, while the lowest among proprietary models, is still a remarkable result for a model with no GPU-specific fine-tuning — it suggests that enough CUDA and HIP code appears in web-scale pretraining data that the model can induce the correspondence rules.

Scaling behavior. As with assembly translation, source accuracy scales monotonically with model size: 1.5B → 90%, 3B → 92.5%, 7B → 95%. The improvement from 1.5B to 7B is +5 percentage points, which is substantial but smaller in relative terms than the assembly scaling (+20 percentage points, 17.5% → 37.5%). This suggests that source translation capability saturates faster with model scale than assembly translation — the CUDA→HIP mapping is shallow enough that even a 1.5B model captures most of it given 70k paired samples, while the SASS→RDNA3 mapping requires larger capacity to continue improving.

Qualitative comparison (Appendix A.5.2). The paper provides three specific examples where CASS-7B outperforms other models in preserving source code semantics:

String constant preservation: CASS-7B retains the original "CUDA" string in printf format strings, while Claude, Qwen-Coder, and GPT-4o all incorrectly change it to "HIP", introducing a semantic error.
Kernel launch syntax: CASS-7B preserves the traditional <<<...>>> launch syntax, while other models replace it with the HIP-specific hipLaunchKernelGGL macro — functionally equivalent but structurally different.
Output stream preservation: CASS-7B retains std::cout as the output stream, while GPT-4o changes it to std::cerr, altering program behavior.

These examples demonstrate that the CASS models learn fidelity to the original semantics, not just syntactic correctness — a property that the automatic accuracy metric captures but qualitative examples make concrete.

Ablation Study: Contribution of Each Data Source and RoPE Extrapolation

Table 4 reports assembly translation accuracy for four ablations of the CASS-7B model:

Data Configuration	Assembly Accuracy (%)
Stack only	17.5
Stack + Synthetic	30.0 (+12.5)
Stack + Synthetic + OpenCL	32.5 (+2.5)
Stack + Synthetic + OpenCL + RoPE	37.5 (+5.0)

Stack data alone (17.5%). The scraped CUDA data from public repositories provides a baseline of 17.5% accuracy. This is the "natural distribution" of GPU code — it represents what real-world CUDA programmers write, as opposed to what LLMs generate when prompted. The fact that Stack data alone achieves 17.5% is evidence that even without synthetic augmentation, the compilation pipeline can extract enough aligned assembly pairs from real-world code to impart non-trivial translation capability. However, 17.5% is far below the full model's 37.5%, indicating that public CUDA repositories have substantial domain and pattern gaps.

Synthetic data (+12.5%). The addition of 46.3k synthetically generated CUDA samples provides the largest single improvement. The paper does not break this down by domain, but the synthetic generation templates (Appendix A.5.1) cover 9 broad categories including simulations, cryptography, optimization, and signal processing — domains that may be underrepresented in public CUDA repositories. The +12.5 percentage point gain suggests that the synthetic data fills coverage gaps, exposing the model to assembly patterns (instruction sequences, memory access patterns, synchronization idioms) that appear in these domains but are rare or absent in Stack data.

OpenCL data (+2.5%). The 6k OpenCL-sourced samples add a smaller but positive improvement. Because OpenCL compiles to both Nvidia and AMD assembly through different paths than CUDA/HIP (OpenCL → PTX → SASS vs. OpenCL → LLVM IR → RDNA3), the assembly pairs it produces have different characteristics — potentially different instruction scheduling patterns, different register allocation strategies, or different handling of synchronization. The +2.5% gain, while modest, is evidence that the model benefits from seeing translation patterns that are not artifacts of the HIPIFY conversion process, consistent with the paper's claim that the CASS dataset captures genuine cross-ISA correspondence rather than HIPIFY-specific mappings.

RoPE extrapolation (+5.0%). The addition of RoPE extrapolation at inference time (extending the 16K-token training context to 32.7K tokens) provides a surprisingly large +5.0 percentage point improvement. This suggests that a non-trivial fraction of assembly files in CASS-Bench exceed the 16K-token training context — without extrapolation, these samples would be truncated, losing portions of the assembly that are necessary for correct translation. The +5.0% gain implies that approximately 5% of samples (2 out of 40) are "rescued" by the extended context, or that the extended context provides partial improvements across multiple samples. This finding has a direct engineering implication: assembly translation is context-length-bound, and larger context windows (either through extrapolation, more efficient tokenization, or architectural changes) are likely to yield accuracy improvements even without additional training data.

What is NOT ablated. The paper does not report ablations on several factors that would strengthen the analysis:

Fine-tuning data volume: How does accuracy scale with the number of training samples? Does doubling the Stack data to 48k (if more repositories were scraped) provide diminishing returns, or is the model still data-hungry at 70k samples?
Model architecture: Would a non-Qwen base model (e.g., DeepSeek-Coder, CodeLlama) achieve similar accuracy with the same data? This is important because Qwen2.5-Coder-32B already achieves 25% assembly accuracy without CASS fine-tuning (Table 3) — perhaps Qwen has unusually good assembly pretraining that other code models lack.
HIPIFY data quality: Since all CUDA-HIP pairs are generated by HIPIFY (with ~44% filtered out), what is the effect of training on HIPIFY-generated translations? Could a model trained on a smaller set of manually verified translations outperform the full CASS model?
Compilation optimization level: The paper uses -Os exclusively. How would the model perform if trained and evaluated on -O0 (no optimization, more verbose but more regular) or -O3 (performance-optimized, more complex mapping)?

Runtime and Memory Fidelity of Translated Assembly

Headline results (Appendix A.4.2, Figure 11). For the assembly translations that compile and execute correctly (the subset corresponding to the 37.5% accuracy figure):

Memory usage: Deviation of less than ±0.3% from native HIP baselines for all files, with 18 files using more memory (maximum +0.3%) and 22 using less (minimum −0.3%).
Execution time: 11 files are slower than native (maximum +11.8%), 8 files are faster (minimum −10.0%), and the rest are unchanged.
Combined metric: "Over 85% of samples fall within ±5.6% across both metrics."

Each test was executed 20 times with "the reported values reflect the average across runs to mitigate noise and ensure statistical reliability" (Appendix A.4.2). This is a standard practice for GPU benchmarking, where runtime variance from system noise, thermal throttling, and driver scheduling can be substantial.

Interpretation of the ±5.6% claim. The paper reports that "over 85% of samples fall within ±5.6% across both metrics" — but this phrasing requires parsing. "Both metrics" means memory usage and execution time. "Within ±5.6%" means that for 85% of correctly translated samples, both the memory usage deviation and the execution time deviation are between −5.6% and +5.6% of the native HIP baseline. This is a strong result: the model not only produces functionally correct assembly, but that assembly has performance characteristics very close to hand-written HIP code (which was originally compiled from the same CUDA source via HIPIFY).

The 11.8% worst-case slowdown. The maximum execution time increase of +11.8% (observed in one file) represents the worst-case performance regression among correctly translated samples. This is substantially larger than the ±5.6% band that captures 85% of samples, indicating that some samples have modest but noticeable performance degradation. The paper does not analyze which domains or which code patterns exhibit the largest slowdowns, which limits the actionable insight — it would be useful to know whether the slowdown occurs in memory-bound kernels (where suboptimal memory coalescing or cache usage could cause regression) or compute-bound kernels (where instruction scheduling inefficiencies could be the culprit).

The -10.0% speedup puzzle. The observation that 8 translated files run faster than the native HIP baseline (by up to 10.0%) is intriguing and unexplained. If the translated RDNA3 assembly is semantically equivalent to the native HIP-compiled RDNA3 assembly, how can it be faster? A few possibilities: (a) the -Os compilation flag used for the native baseline produces conservative code, and the model generates slightly more aggressive instruction sequences that happen to run faster; (b) measurement noise — despite 20-run averaging, GPU runtime variance can be large for short-running kernels (if the benchmark includes kernels that run in microseconds, even 20-run averaging may not fully stabilize); (c) the model's translations happen to trigger different compiler optimization paths in the AMD driver's JIT compilation of the generated assembly, leading to different final instruction schedules. The paper does not investigate which of these explanations applies.

What this measurement establishes (and does not establish). The memory and runtime measurements establish that the CASS model's translations, when correct, preserve the computational structure of the original code. This is evidence against the hypothesis that the model is merely producing "plausible-looking" assembly that happens to produce the correct output through a very different computational path — if that were the case, runtime and memory characteristics would diverge substantially. Conversely, the measurements do not establish that the model can reproduce hardware-specific optimizations from the SASS source — the fact that some translated files run up to 11.8% slower than native HIP code suggests that some low-level optimizations are being lost in translation. This is consistent with the paper's candid admission that the model is "not optimization-aware yet" (Section 1).

Critical Assessment

The experiments in this paper demonstrate something genuine and important: that a language model fine-tuned on aligned, paired GPU assembly data can learn to translate between NVIDIA SASS and AMD RDNA3 assembly at a level that enables functional correctness for a non-trivial fraction of test cases, while all general-purpose models — including GPT-4o and Claude-3.7 — score 0% on the same task. This is a valid and well-supported claim: Table 3 shows CASS-7B at 37.5% vs. 0% for all commercial baselines, and the 40-sample CASS-Bench, while small, is manually verified for output equivalence, making the measurement trustworthy.

However, several aspects of the experimental design limit how broadly we should interpret these results, and some of the paper's implicit claims are not adequately tested.

The 37.5% assembly accuracy is a meaningful but fragile number. With 40 test samples and 37.5% accuracy, the 7B model correctly translates exactly 15 out of 40 CASS-Bench samples. A single additional correct or incorrect translation would shift the reported accuracy by 2.5 percentage points. The difference between CASS-7B (37.5%) and CASS-3B (25.0%) — 5 additional correct translations — is suggestive of scaling but not statistically robust. The paper reports no confidence intervals or significance tests. Moreover, the domain breakdown (0% in math/data structures/graph, 25–50% in linear algebra, 100% in physics) suggests that the 37.5% average is heavily influenced by which domains happen to be represented in CASS-Bench. A CASS-Bench with different domain weighting (e.g., more math problems, fewer physics simulations) would yield substantially different average accuracy. The paper's headline number should therefore be understood as a demonstration of feasibility on a specific benchmark distribution, not as a reliable estimate of expected accuracy on arbitrary GPU workloads.

The paper demonstrates learning but does not establish generalization. The CASS models are trained on data that was generated through a specific pipeline: CUDA → HIPIFY → HIP → compilation → assembly extraction. CASS-Bench samples are constructed through the same pipeline — the benchmark's AMD reference code was generated by Claude-3.7 and then verified, but the assembly pairs were extracted using the same compilation infrastructure (Section 4.2). This means CASS-Bench evaluates in-distribution generalization: can the model handle new samples from the same data-generating process? This is a valid evaluation, but it leaves open the critical question of out-of-distribution generalization: can the model translate CUDA code that was NOT first converted through HIPIFY, or handle SASS from GPU architectures other than sm_85, or translate RDNA3 from compilation configurations other than -Os? The paper does not test any of these generalization scenarios. The finding that OpenCL data (which does NOT go through HIPIFY) contributes +2.5% to accuracy (Table 4) provides weak positive evidence for generalization, but this is an ablation on training data composition, not an evaluation of generalization to unseen data distributions.

The HIPIFY comparison is both a strength and a confound. The paper's source translation results show CASS-7B outperforming HIPIFY by 7.5 percentage points (95% vs. 87.5% on CASS-Bench). This is a genuine achievement — the model learned to handle edge cases that HIPIFY's rule-based translation misses. However, HIPIFY's 87.5% on CASS-Bench is not representative of its performance on arbitrary CUDA code, because the CASS-Bench samples are "HIPIFY-compatible" by construction (they were generated by Claude-3.7, compiled, and passed through the CASS pipeline, which filters out HIPIFY failures). The paper's reported 43.9% HIPIFY failure rate on the training data pipeline suggests that on a random sample of real-world CUDA code, HIPIFY would score substantially below 87.5%. The CASS model might score substantially below 95% on the same random sample. The 7.5% improvement should therefore be interpreted as improvement on the subset of CUDA code that HIPIFY can handle at all, not as the improvement on arbitrary CUDA.

The 0% accuracy of GPT-4o and Claude-3.7 requires methodological scrutiny. Table 3 reports 0% for both models on assembly translation. The paper does not describe the prompting strategy used for these baselines — what instructions were given, whether few-shot examples were provided, what output format was requested, whether the models were asked to output raw assembly or were allowed to explain their reasoning. This matters because assembly translation is a highly specific task that general-purpose models may not understand from a naive prompt. If GPT-4o was simply asked "Translate this SASS to RDNA3" without context, examples, or format specification, its 0% accuracy may reflect task misunderstanding rather than fundamental incapability. The paper's note that Qwen2.5-Coder-32B achieved 25% "when explicitly prompted" (Table 3, footnote) suggests that prompting matters. But the paper does not apply the same "explicit prompting" methodology to GPT-4o or Claude-3.7, making the 0% comparison potentially unfair. A rigorous comparison would have included best-effort prompt engineering for each baseline model, perhaps including few-shot examples from the training set.

The missing baseline: what would a simple n-gram or SMT model achieve? The paper frames its contribution as demonstrating that "domain-specific fine-tuning on aligned cross-vendor assembly pairs can impart low-level ISA translation knowledge." But 37.5% accuracy leaves open the question of how much of this capability comes from the neural architecture and how much comes from the data itself. A phrase-based statistical machine translation model or even a sophisticated n-gram alignment model, trained on the same 70k paired samples, might achieve non-zero accuracy — perhaps 15-20%. If so, the neural model's contribution would be more modest than the "0% to 37.5%" framing suggests. The paper includes no non-neural baselines (IBM Model 1, Moses, etc.) that would calibrate expectations for what the data alone enables.

Performance fidelity measurements need more analysis to be fully convincing. The finding that 85% of translated samples fall within ±5.6% of native performance for both memory and runtime (Appendix A.4.2, Figure 11) is reported as supporting evidence for the model's quality. But these measurements apply only to the correctly translated subset — the 15 out of 40 samples (at 37.5% accuracy) that passed functional verification. We do not know the performance characteristics of the incorrectly translated samples (do they crash? produce wrong results with plausible-looking assembly? produce assembly that is correct but with wildly different performance?). Moreover, the performance measurements compare the model's RDNA3 output against the native HIP baseline, not against the NVIDIA SASS source. The original motivation for assembly-level translation was to preserve hardware-specific optimizations — but these measurements cannot tell us whether optimizations were preserved because they compare against what the AMD compiler produces from HIP source, not against what the NVIDIA compiler embedded in SASS. A more informative comparison would measure whether the translated RDNA3 assembly achieves speedups relative to HIP-compiled RDNA3 on kernels where the NVIDIA SASS source contains hand-tuned optimizations — but CASS-Bench does not contain such kernels because it is constructed from LLM-generated code, not hand-optimized CUDA.

The missing experiment: training on HIPIFY failures. The paper reports that 43.9% of CUDA files fail HIPIFY conversion and are discarded. These failures represent the hardest cases — CUDA code using unsupported features, complex templates, or non-standard patterns. Training a model on these failures (with manually written HIP translations) and testing whether the model can handle unsupported CUDA features would be a much stronger demonstration of the value of learned translation over rule-based approaches. The paper does not attempt this, and as a result, the CASS models are trained and evaluated entirely on HIPIFY-compatible code — they have not been tested on the very cases that motivate the need for ML-based translation in the first place.

The 1.5B model's 90% source accuracy raises questions about data efficiency. Even the smallest CASS model (1.5B parameters, trained on 70k samples) achieves 90% source accuracy, outperforming GPT-4o (80%) and Claude-3.7 (85%). This suggests that source-level CUDA→HIP translation is a relatively easy task that saturates quickly with data — the model essentially needs to learn a finite set of API-call substitutions and a few templating conventions. The paper does not explore whether even smaller datasets (10k, 1k, or 100 samples) would achieve comparable accuracy, which would inform whether the full 70k dataset is necessary for source translation or whether the dataset's primary value is for assembly translation.

In summary: The experiments convincingly demonstrate that (a) general-purpose models have essentially zero GPU assembly translation capability, (b) fine-tuning on the CASS dataset imparts non-trivial capability (37.5% accuracy), and (c) the CASS pipeline generates training data that enables this learning. These are valid and important demonstrations. The experiments do NOT establish that the learned capability generalizes beyond the specific compilation pipeline, GPU architectures, and code patterns present in CASS-Bench; that the models can translate code that HIPIFY cannot; that prompt engineering cannot extract assembly translation capability from larger general-purpose models; or that the CASS dataset size is necessary rather than merely sufficient for source translation. These gaps do not invalidate the paper's contribution — which is infrastructural — but they bound the claims that can reasonably be made about the trained models' practical utility.

6. Limitations and Trade-offs

6.1 Assembly Translation Accuracy Is Domain-Dependent and Zero for Several Important GPU Workload Categories

The assumption or constraint. The paper implicitly assumes that a 70k-sample dataset covering 16 GPU domains will provide sufficient coverage for assembly translation to be useful across common GPU workloads. However, the reported 37.5% average accuracy on CASS-Bench masks extreme heterogeneity: the model achieves 0% accuracy on math, data structures, and graph algorithm tasks, while reaching 100% on physics simulations (Section 6, Figure 5; Appendix A.4.2). The paper acknowledges this in passing: "assembly accuracy is inconsistent, 0% in Math, Data Structures, and Graph, 25–50% in linear algebra and memory operations" (Appendix A.4.2), but does not analyze why these domains fail or what fraction of real-world GPU workloads they represent.

The consequence. A practitioner cannot rely on CASS assembly translation for programs containing irregular memory access patterns (pointer chasing, sparse data structures), complex control flow (graph traversals, recursive algorithms), or mathematical kernels not well-represented in the training data. These failure modes align with computational patterns where the mapping between SASS and RDNA3 instruction sequences is most context-dependent — register allocation, memory addressing modes, and branch prediction interact in ways that a sequence-to-sequence model apparently cannot capture from 70k examples, or cannot capture because these domains are underrepresented in the training set. The paper's domain distribution analysis (Figure 3, left) shows training data coverage of these categories, but the specific sample counts per domain are not reported, making it impossible to distinguish between "insufficient training data" and "fundamentally harder translation problem" as the root cause.

What evidence exists in the paper. Figure 5 (right panel) shows the per-category assembly accuracy breakdown, with math, data structures, and graph tasks at 0%. Figure 3 (right) shows CASS-Bench contains samples from these domains (1–5 per category). Table 4 shows that adding synthetic data (+12.5%) and OpenCL data (+2.5%) improves overall assembly accuracy, but the ablation does not report per-domain improvements — it is unknown whether these data sources help specifically in the 0%-accuracy domains. The paper does not provide per-domain training sample counts, making it impossible to assess whether the 0% accuracy categories are data-sparse or whether they receive adequate training coverage but remain unlearnable.

Mitigation status. The paper does not attempt to address domain-specific failures. Section 7 (Limitations and Future Work) acknowledges that "current performance is inadequate for production due to limited accuracy in complex or underrepresented domains" and that "expanding category diversity is essential," but this is a general statement rather than a targeted mitigation plan for the 0%-accuracy categories. No analysis is provided of what makes math, data structure, and graph algorithm assembly harder to translate than physics simulation assembly, so practitioners have no guidance on whether collecting more data in these domains would help (suggesting a data volume problem) or whether architectural changes are needed (suggesting a model capacity or inductive bias problem).

6.2 The Dataset Covers Only One GPU Architecture Pair Per Vendor (sm_85 ↔ RDNA3 on RX 7900 XT)

The assumption or constraint. All training and evaluation data targets exactly one Nvidia ISA (sm_85, the Ampere architecture SASS produced by an A100 GPU) and one AMD ISA (RDNA3, produced by an RX 7900 XT GPU). The paper explicitly acknowledges this: "The dataset currently covers only one host/device pair per vendor (RTX 4090 and RX7900), limiting generalizability across GPU architectures with varying ISAs" (Section 7). Additionally, although an A100 was used for compilation, the CUDA code was compiled targeting consumer-grade compute capabilities (RTX 4090) "to maintain parity with the AMD hardware" (Appendix A.2) — meaning the SASS instructions exclude A100-specific features like certain tensor core operations.

The consequence. The trained CASS models have no exposure to instruction set differences across GPU generations from the same vendor. Nvidia's SASS evolves substantially between architectures — Volta (sm_70), Turing (sm_75), Ampere (sm_80/sm_86), Ada Lovelace (sm_89), and Hopper (sm_90) each introduce new instructions, deprecate old ones, and change performance characteristics of existing instructions. AMD's RDNA1, RDNA2, and RDNA3 ISAs similarly diverge. A model trained only on sm_85 ↔ RDNA3 pairs has no knowledge of how to map a Hopper-specific STMATRIX instruction or an RDNA2-specific v_pk_add_f16 instruction. For practitioners with codebases targeting different GPU generations (which is nearly all real-world GPU software, since applications typically support multiple hardware generations), the current CASS models would fail on any instruction outside the sm_85/RDNA3 intersection. The paper also notes ZLUDA's 2.5% accuracy may reflect "its compatibility with RDNA1" (Section 6) — implying that even a single generation gap (RDNA1 → RDNA3) causes substantial failure, and by analogy, the CASS models would similarly fail on non-Ampere SASS or non-RDNA3 targets.

What evidence exists in the paper. The limitation is acknowledged explicitly in Section 7, but no experiment quantifies the magnitude of the generalization gap. The paper does not evaluate CASS models on SASS from a different Nvidia architecture (e.g., Turing sm_75) or on RDNA2 assembly, so the cross-generation failure rate is unknown. Table 3 shows that the base Qwen2.5-Coder-32B achieves 25% assembly accuracy without CASS fine-tuning — this could reflect the base model having seen SASS and RDNA3 from multiple architectures during pretraining (since The Stack includes code targeting various GPU generations), but the paper does not analyze whether this 25% comes from sm_85-specific patterns or from architecture-agnostic patterns that would transfer.

Mitigation status. Section 7 states that "broader architectural representation is needed to support real-world deployment" but provides no concrete plan for how to extend the CASS pipeline to additional GPU architectures. The pipeline is designed to be architecture-parameterized (one specifies the compilation target when invoking nvcc or hipcc), so adding new architectures is an engineering effort rather than a methodological challenge — but it requires access to the additional GPU hardware for compilation and verification, which is a non-trivial resource barrier. The paper does not discuss whether the existing pipeline can generate multi-architecture data without hardware access (e.g., cross-compilation), or whether models trained on sm_85/RDNA3 could be fine-tuned on smaller datasets for additional architectures.

6.3 The Difficulty Estimation and Filtering Pipeline Introduces Selection Bias Favoring HIPIFY-Compatible and Compilable Code

The assumption or constraint. The CASS dataset construction pipeline systematically excludes two categories of CUDA code: (1) code that HIPIFY cannot convert (~43.9% of CUDA files, Section 3.3), and (2) code that fails compilation on either the Nvidia or AMD stack (discards asymmetric failures, Section 3.3). The paper acknowledges the HIPIFY failure rate but does not discuss its implications for the trained models' capabilities. The compilation filter additionally removes all synthetic samples with syntactic errors, missing definitions, or invalid memory operations (the 50.9% of generated CUDA that does not compile; Section 3.2).

The consequence. The CASS models are trained exclusively on HIPIFY-compatible CUDA code that compiles cleanly under both Nvidia and AMD toolchains. This means the models have never seen (and cannot learn to translate) CUDA patterns that cause HIPIFY failures — unsupported API calls, complex template metaprogramming, features from recent CUDA versions that HIPIFY has not yet adopted. These are precisely the cases where a learned translator would be most valuable, since HIPIFY already handles the straightforward cases. The models also have no exposure to CUDA code with subtle bugs (off-by-one errors, race conditions, uninitialized memory) that survive compilation — such code exists in real codebases but is filtered out because it fails the pipeline's execution verification or produces non-matching outputs. The resulting models may be brittle when encountering real-world CUDA that is not "pipeline-clean."

Similarly, the synthetic data filtering (removing 50.9% of generated samples that fail compilation) means the models only see LLM-generated CUDA that is already compilable — they never learn to handle or correct the kinds of errors that LLMs commonly make when generating GPU code. This limits the models' usefulness in a pipeline where the input CUDA is itself LLM-generated (an increasingly common scenario).

What evidence exists in the paper. The 43.9% HIPIFY failure rate is reported in Section 3.3. The 49.1% synthetic compilation success rate (i.e., 50.9% failure rate) is reported in Section 3.2. The paper does not report the compilation failure rate for Stack-scraped CUDA (the 24k usable samples are post-filtering, so the raw compilation rate is unknown). Table 3 shows the trained models outperform HIPIFY by 7.5% on CASS-Bench, but CASS-Bench is itself constructed from HIPIFY-compatible, compilation-verified code (Section 4.2) — so this 7.5% improvement is measured on the selected subset that HIPIFY can already handle, not on the 43.9% of code that HIPIFY fails on entirely. The paper does not evaluate CASS models on HIPIFY-failure cases, so the models' capability on truly hard translation cases is unknown.

Mitigation status. Not addressed. The paper presents the filtering steps as quality-control measures (removing uncompilable or unconvertible code to produce a clean dataset) without discussing the resulting distribution shift between the training data and real-world CUDA code. Section 7 acknowledges that "dataset size was minimized to fit within 16K-token context windows, excluding many vendor-specific low-level optimizations" — this is a related but distinct limitation (context window constraints) rather than a discussion of selection bias. No experiment tests model performance on HIPIFY-failure cases, and no future work is proposed to address the coverage gap (e.g., manually translating a sample of HIPIFY-failure cases to create an out-of-distribution test set).

6.4 The CASS-Bench Evaluation Set (40 Samples) Is Too Small and Too Homogeneous to Support Robust Quantitative Claims

The assumption or constraint. All model comparisons and accuracy claims rest on a 40-sample evaluation set spanning 16 domains with 1–5 samples per domain (Section 4.2, Figure 3 right). The paper treats per-domain accuracy as meaningful even for domains with only 1–2 samples, and reports overall accuracy averages (37.5%, 95%) without confidence intervals, standard deviations, or significance tests. The 40 CASS-Bench samples are all generated by Claude-3.7 from natural language prompts and manually verified for output equivalence — they share a common "generation style" (the idioms, coding patterns, and optimization choices of Claude-3.7) that may not represent the diversity of human-written CUDA code.

The consequence. The headline accuracy numbers are fragile. At 37.5% assembly accuracy, the CASS-7B model correctly translates exactly 15 out of 40 samples. A single additional correct or incorrect prediction shifts the reported accuracy by 2.5 percentage points. The difference between CASS-7B (37.5%, 15/40 correct) and CASS-3B (25.0%, 10/40 correct) — a 12.5 percentage point gap representing 5 additional correct translations — is suggestive of model scaling benefits but could arise from variance in a small sample. Per-domain accuracies with 1–2 samples are essentially meaningless as quantitative measurements: a domain with 1 sample and 100% accuracy tells us nothing about whether the model would succeed on a second sample from the same domain. The 100% accuracy on physics simulation (Appendix A.4.2) could represent 1/1, 2/2, or up to 5/5 correct — the paper does not report per-domain sample counts in CASS-Bench.

The homogeneity concern compounds the small-sample problem. All 40 samples were generated by the same model (Claude-3.7) using similar prompting methodology. Claude-3.7 has consistent coding style — particular preferences for variable naming, loop structure, error handling, and optimization patterns. If CASS models learn to translate "Claude-3.7-style CUDA" but fail on human-written CUDA with different conventions, the CASS-Bench evaluation would overestimate real-world performance. The paper provides qualitative examples (Appendix A.5.2) where CASS-7B correctly handles edge cases (preserving string constants, kernel launch syntax) that other models mishandle, but these examples are for source translation and do not address whether Claude-generated assembly has systematic differences from human-written assembly.

What evidence exists in the paper. Table 3 reports all accuracy numbers without error estimates. Figure 5 shows per-domain breakdowns without sample counts. Figure 10 (Appendix) shows accuracy vs. training steps curves for the three model scales, but the curves appear smooth (suggesting training stability, not that the test set is large enough for precise measurement). The paper acknowledges in Section 7 that "current performance is inadequate for production due to limited accuracy," but does not identify the evaluation set size as a limitation. No experiment tests CASS model performance on human-written CUDA not generated by Claude-3.7 (e.g., by evaluating on a held-out subset of Stack-scraped CUDA with manually written HIP translations).

Mitigation status. Not addressed. The paper presents CASS-Bench as a contribution ("the first evaluation benchmark for cross-architecture GPU translation") without discussing its statistical limitations. Section 4.2 describes the curation process in detail but does not justify the 40-sample size or discuss power calculations. No future work is proposed to expand CASS-Bench or to develop statistical evaluation protocols (e.g., bootstrap confidence intervals, significance testing for model comparisons). A larger CASS-Bench with human-written CUDA samples and per-domain sample counts large enough for statistical analysis is a clear next step, but the paper does not flag it as such.

6.5 The Baseline Comparisons Are Not Rigorously Controlled, Particularly for Assembly Translation

The assumption or constraint. The paper's most striking claim — that CASS models achieve 37.5% assembly accuracy while GPT-4o, Claude-3.7, and Gemini 2.0 Flash all score 0% (Table 3) — depends on the assumption that the baseline models were given a fair opportunity to succeed at the task. The paper does not describe the prompting methodology used for any baseline model (Section 6), beyond a footnote that Qwen2.5-Coder-32B achieved its 25% accuracy "when explicitly prompted" (implying that prompting matters and that the 0% models may have been evaluated with suboptimal prompts). The baselines are evaluated on a task — raw SASS-to-RDNA3 translation — that these models were never trained for and may not even recognize as a valid request without careful instruction.

The consequence. The 0% results for proprietary models may reflect task framing failure (the model did not understand what was being asked, or refused the task, or produced output in an unparseable format) rather than fundamental incapability. A general-purpose model like GPT-4o, if provided with few-shot examples of SASS → RDNA3 translation pairs, a detailed system prompt explaining the ISA correspondence, and explicit output formatting instructions, might achieve non-zero accuracy — perhaps 5–15%, which would substantially change the narrative from "domain-specific training is essential" to "domain-specific training provides substantial improvement over prompted general models." The paper provides no evidence to distinguish these alternatives.

The Qwen2.5-Coder-32B result (25% "when explicitly prompted") demonstrates that a model without CASS fine-tuning can achieve non-trivial assembly accuracy given appropriate prompting. Since Qwen2.5-Coder-32B is architecturally similar to the 7B base model used for CASS fine-tuning, this 25% represents an upper bound on what prompting alone can achieve — and it is a substantial fraction of the 37.5% achieved after fine-tuning on 70k samples. This narrows the contribution of the CASS dataset: fine-tuning adds +12.5 percentage points over the best prompted baseline (25% → 37.5%), not the +37.5 points that the "0% vs. 37.5%" framing implies. If GPT-4o or Claude-3.7 could achieve even 10–15% with optimized prompting, the CASS dataset's marginal contribution would appear even more modest.

What evidence exists in the paper. Table 3 reports baseline accuracies without describing prompts, number of attempts, output parsing methodology, or error analysis. The Qwen2.5-Coder-7B-Instruct result (0%) is particularly puzzling — the instruction-tuned variant of the same model family scores 0% while the base model scores 25% — and the paper provides no explanation beyond the "explicitly prompted" footnote. This suggests that instruction tuning on general code tasks may actively suppress assembly translation capability (perhaps by training the model to refuse requests it considers unusual or to produce explanatory text rather than raw assembly), but this hypothesis is not investigated. The paper does not report whether the 0%-scoring models produced any output at all, produced plausible-looking assembly that failed compilation, or refused the task entirely.

Mitigation status. Not addressed. The paper does not describe any effort to optimize prompts for baseline models, does not provide few-shot examples (which the CASS dataset could easily provide, since the training set contains 70k aligned pairs), and does not analyze failure modes of the baseline models. Section 7 does not identify baseline fairness as a limitation. A rigorous follow-up would include: standardized prompt templates with task descriptions, few-shot examples drawn from CASS-Instruct (not CASS-Bench), multiple prompt variants per model, and reporting of both best-of-N and average accuracy across prompts. None of this is done in the current paper.

6.6 Translation Latency and Hardware Requirements Make the Assembly Translation Pipeline Impractical for Interactive or Real-Time Use

The assumption or constraint. The paper evaluates the CASS models in terms of accuracy and runtime fidelity of the generated code, but does not account for the end-to-end latency of the translation process. The inference cost is reported as "approximately 56 seconds per a 16K-token sample" on 4× A100 GPUs (Section 5). This is the time to generate a single translated assembly file — it does not include the time to compile and verify the translation, or to iterate if the translation fails compilation (which would happen for 62.5% of assembly translation attempts at the current 37.5% accuracy).

The consequence. A 56-second per-translation latency, requiring 4 data-center-class GPUs (A100s), makes assembly translation completely impractical for interactive development workflows (where a developer expects near-instant feedback), for just-in-time compilation scenarios (where translation must happen at application launch time), or for batch processing of large codebases (where translating thousands of kernels would take days of GPU time). This is not a minor engineering concern — it fundamentally changes the use cases for which assembly translation is viable. The only plausible deployment scenario given current latency is offline, batch translation of critical kernels where the translation cost can be amortized over many subsequent executions, and where the alternative (manual HIP rewrite) would be even more expensive.

The hardware requirement (4× A100 GPUs for the 7B model) further restricts deployment to organizations with substantial GPU infrastructure — the same organizations that could afford to maintain both Nvidia and AMD hardware and might have less need for cross-vendor translation. Smaller organizations or individual developers who would most benefit from automated CUDA-to-AMD migration cannot run the 7B model locally. The 1.5B model (17.5% assembly accuracy) might run on consumer hardware, but the paper does not report its inference latency or hardware requirements.

What evidence exists in the paper. The 56-second figure is reported in Section 5 without further analysis. The 4× A100 hardware configuration is described in the same section. The paper does not report inference latency for the 1.5B or 3B models, does not discuss whether model quantization or distillation could reduce latency, and does not frame latency as a limitation. Section 7 mentions that "dataset size was minimized to fit within 16K-token context windows" but does not connect this to inference latency or discuss whether smaller models with lower latency could achieve acceptable accuracy.

Mitigation status. Not addressed as a limitation. The paper treats inference as efficient ("Inference was efficient, requiring approximately 56 seconds per a 16K-token sample," Section 5) without contextualizing what "efficient" means for practical use cases. The 98% GPU utilization during training (Section 5) is reported as an achievement, but inference efficiency (tokens per second, latency per sample, GPU memory requirements per model scale) is not analyzed. The paper does not propose model compression, quantization, speculative decoding, or other latency-reduction techniques as future work, focusing instead on accuracy improvements and domain coverage expansion.

7. Implications and Future Directions

How This Work Changes the Landscape

This paper shifts the field of GPU cross-architecture translation from a compiler engineering problem (build a better transpiler) to a data infrastructure and machine learning problem (build a paired dataset, then train a model). This is not merely a semantic reframing—it has concrete consequences for what kinds of work become tractable and what kinds of solutions the community will pursue.

Before CASS, the dominant approaches to CUDA-to-AMD migration—HIPIFY, CuPBoP-AMD, ZLUDA, GPU Ocelot—were all rule-based or IR-level systems built by compiler engineers who manually encoded translation rules. Each of these tools hit fundamental ceilings: HIPIFY fails on ~44% of CUDA files (Section 3.3), ZLUDA achieves only 2.5% assembly accuracy on CASS-Bench (Table 3), and GPU Ocelot was abandoned due to scalability issues. The implicit assumption was that these failures reflected algorithmic difficulty—that GPU ISAs are too divergent for simple translation rules and that we needed more sophisticated compiler techniques.

CASS demonstrates that this diagnosis is incomplete. It shows that a relatively straightforward sequence-to-sequence model (fine-tuned Qwen2.5-Coder with standard supervised learning) can achieve non-trivial assembly translation accuracy (37.5% on CASS-Bench) when provided with 70k aligned, paired assembly examples—a resource that simply did not exist before. The fact that all general-purpose models score 0% on the same task (GPT-4o, Claude-3.7, Gemini 2.0 Flash; Table 3) confirms that the missing ingredient was data, not algorithmic sophistication. This reframing has several downstream effects on the research landscape:

First, it makes GPU assembly translation a data problem, not a compiler problem. The ablation results in Table 4 tell a clear story: Stack data alone → 17.5% accuracy; add synthetic data → 30% (+12.5); add OpenCL data → 32.5% (+2.5); add RoPE extrapolation → 37.5% (+5.0). Every data contribution improves the model, with no change to the translation algorithm. This suggests that the marginal return to better data is currently high, and that future work should invest in expanding and diversifying the dataset rather than designing cleverer model architectures. The paper makes the CPU assembly translation precedent (CRT, Guess & Sketch) newly applicable to GPUs by providing the paired corpus that those approaches depend on but that never existed for GPU ISAs.

Second, it establishes a new diagnostic for GPU ISA compatibility. The finding that assembly accuracy varies from 0% (math, data structures, graph algorithms) to 100% (physics simulations) reveals that the SASS→RDNA3 mapping is not uniformly learnable—its difficulty depends on the computational pattern being expressed. This was not predictable a priori. It tells us that regular, floating-point-heavy computations produce systematically translatable assembly while irregular, control-flow-heavy computations do not. This is a structural discovery about the relationship between GPU ISAs that can guide future work: researchers now know which domains to prioritize for data collection (the 0% categories) and which domains can serve as easier testbeds for methodological development (the 100% categories).

Third, it reconciles a latent contradiction in the literature. The prior work on GPU portability presented a confusing picture: HIPIFY works for simple CUDA but fails on ~44% of real code; ZLUDA promises binary translation but achieves near-zero accuracy on modern GPUs; CPU assembly translation succeeds with language models but no one had tried the approach for GPUs. CASS provides a unified explanation: the GPU case was blocked by a data availability problem, not a fundamental incompatibility between GPU ISAs. The CPU success stories (CRT, Guess & Sketch) were not algorithmically novel—they simply had access to paired corpora that could be constructed from compiler output. CASS fills the analogous gap for GPUs, making the CPU precedent actionable. This resolution is valuable because it converts a set of seemingly contradictory results (HIPIFY works partially, ZLUDA works poorly, language models work not at all) into a coherent picture where the common missing factor is aligned training data.

Fourth, it introduces execution fidelity as a co-emergent property worth measuring. Before CASS, GPU translation evaluation focused on compilation success or output correctness. The paper's finding that over 85% of correctly translated samples preserve runtime and memory within ±5.6% of native HIP baselines (Appendix A.4.2, Figure 11) establishes that token-level accuracy on a paired-assembly dataset correlates with performance preservation—the model is not just producing plausible-looking assembly that happens to compute the right answer through a wildly different computational path. This means future work can evaluate on CASS-Bench accuracy and have reasonable confidence that accuracy improvements will translate to usable code, dramatically lowering the evaluation barrier for new approaches.

One thing this work does NOT change: it does not make assembly translation production-ready. The 37.5% accuracy is on a small, Claude-generated benchmark with in-distribution evaluation. The paper is candid about this, and the landscape shift is in making the problem tractable, not in solving it. Researchers who previously would have dismissed GPU assembly translation as too hard or too data-sparse now have a concrete starting point, a benchmark, and a baseline model—and that is the meaningful change.

Follow-Up Research This Work Enables

Out-of-distribution evaluation on real-world, human-written CUDA codebases. The CASS-Bench evaluation set comprises 40 samples all generated by Claude-3.7 from natural language prompts (Section 4.2), meaning they share Claude's coding style, optimization patterns, and idiomatic conventions. A critical follow-up is to construct an evaluation set from human-written CUDA code—drawing from the same Stack v2 repositories that supplied the training data but selecting code that was NOT used in training—and manually produce verified HIP translations for these samples. This would answer: does the 37.5% assembly accuracy hold for real-world CUDA that uses idioms, patterns, and coding styles different from Claude's? The ablation in Table 4 shows that Stack-only training data achieves 17.5% (vs. 30% with synthetic data), suggesting the model benefits from the synthetic data's coverage—but we don't know if the model has overfit to the specific generation patterns of Qwen2.5-Coder 32B (the synthetic data generator) and Claude-3.7 (the benchmark generator). A test set of 50–100 human-written CUDA kernels from diverse open-source projects (e.g., TensorFlow custom ops, GROMACS molecular dynamics, OpenCV GPU modules) with manually verified HIP translations would stress-test whether the learned SASS→RDNA3 mapping is general or generation-style-specific.

Performance-aware translation with optimization-level conditioning. The paper uses -Os (optimize for size) exclusively, achieving 9.3% token reduction relative to -O3 (Section 3.3). This choice was driven by the 16K-token context window constraint, not by a belief that -Os produces optimal translations. A natural extension is to train separate CASS models on -O0, -O2, and -O3 compiled assembly pairs, potentially with an optimization-level token prepended to the input (e.g., <OPT=O3>), enabling the model to generate assembly at different optimization levels on demand. The key experiment: for a set of compute-bound kernels where -O3 significantly outperforms -Os on the native HIP baseline, does the model trained on -O3 pairs produce translated RDNA3 assembly with runtime closer to the -O3 native baseline than a model trained on -Os pairs? If so, this establishes that optimization-level information can be transmitted through the translation, addressing a major limitation of the current approach (the paper explicitly states the model is "not optimization-aware yet," Section 1). A negative result—that the model cannot learn optimization-level-specific mappings—would suggest that the SASS→RDNA3 mapping at -O3 is too complex for the current architecture and that performance-aware translation requires fundamentally different approaches (e.g., searching over multiple generated translations and selecting the fastest).

Cross-generational GPU ISA translation via multi-architecture training. All CASS training data targets exactly one Nvidia ISA (sm_85, Ampere) and one AMD ISA (RDNA3 on RX 7900 XT). The paper acknowledges this as a limitation (Section 7). A high-impact extension would be to extend the CASS pipeline to generate aligned pairs for additional GPU architectures—for Nvidia: sm_75 (Turing), sm_86 (Ampere with RT cores), sm_89 (Ada Lovelace), sm_90 (Hopper); for AMD: RDNA1, RDNA2. The question is whether a model trained on multi-architecture pairs can translate between ISAs it has never seen paired—for instance, training on sm_85↔RDNA3 and sm_75↔RDNA3 but testing on sm_85↔sm_75 (same-vendor cross-generation translation). If this works, it would demonstrate that the model is learning an architecture-agnostic representation of GPU computation, not just memorizing instruction correspondences for specific ISA pairs. The experiment requires access to multiple GPU generations from each vendor—a resource constraint, but one that is feasible for well-equipped academic or industry labs. A negative result (no cross-generational transfer) would indicate that the learned mapping is surface-level and that each architecture pair needs its own training data—a practically important finding that would shape how the community invests in dataset construction.

Prompting baselines with few-shot examples from CASS-Instruct to establish a fair comparison. The paper reports 0% assembly accuracy for GPT-4o, Claude-3.7, and Gemini 2.0 Flash (Table 3) but provides no detail on prompting strategy. The footnote that Qwen2.5-Coder-32B achieved 25% "when explicitly prompted" suggests prompting matters substantially. A rigorous follow-up would provide each commercial model with: a detailed task description, 3–5 few-shot SASS→RDNA3 translation examples drawn from CASS-Instruct (not CASS-Bench), explicit output formatting instructions, and possibly chain-of-thought prompting (explain the translation before outputting it). The question is not "can these models beat CASS-7B?" (they almost certainly cannot) but "what is the true zero-shot and few-shot capability ceiling for general-purpose models on this task?" If GPT-4o or Claude-3.7 can achieve 10–20% with optimized prompting, the marginal contribution of the CASS dataset shifts from "37.5% vs. 0%" to "37.5% vs. 15%," which is still substantial but provides a more honest baseline. If they remain at 0% despite best-effort prompting, that is an even stronger result for the paper's central claim that general-purpose pretraining does not impart GPU assembly translation capability. The key methodological requirement: publish the exact prompts, few-shot examples, and evaluation protocol so the comparison is reproducible.

Structured prediction with architectural constraints for the 0%-accuracy domains. The paper finds 0% assembly accuracy on math, data structures, and graph algorithm tasks (Section 6, Figure 5). These domains likely fail because they involve irregular control flow, pointer chasing, and complex register allocation—patterns where the SASS→RDNA3 mapping is too context-dependent for a sequence-to-sequence model to capture from 70k examples. A targeted follow-up would explore whether adding architectural constraints—such as a constraint solver that ensures register liveness consistency between SASS and RDNA3, or a symbolic execution engine that verifies control flow equivalence—can rescue performance in these domains. The experimental design: take the 0%-accuracy CASS-Bench samples, provide the model with the SASS input and allow it to generate candidate RDNA3 translations, then use a verifier (e.g., a lightweight symbolic execution tool that checks whether the RDNA3 control flow graph matches the SASS control flow graph) to filter or rerank candidates. The question is whether the primary failure mode in these domains is the model producing plausible-looking but semantically incorrect assembly (which a verifier might catch), or the model producing assembly that fails to compile entirely (which a verifier cannot help with). The Guess & Sketch approach (Lee et al., 2023, reference [25]) provides a template for this integration, and the CASS dataset provides the training data to make the language model component viable for GPUs.

Scaling laws for GPU assembly translation data volume. The paper trains on 70k samples and shows scaling with model size (1.5B → 17.5%, 3B → 25%, 7B → 37.5%), but does not explore scaling with data volume. A systematic study would train the 7B model on random subsets of CASS-Instruct at sizes 1k, 5k, 10k, 25k, 50k, and the full 70k, measuring assembly accuracy on CASS-Bench at each data volume. This would reveal whether the model is data-saturated at 70k (suggesting architectural improvements are needed for further gains) or still improving (suggesting that expanding the dataset—e.g., to 200k or 500k samples—would yield substantial accuracy improvements). The ablation in Table 4 provides partial evidence (Stack → Stack+Synthetic → Stack+Synthetic+OpenCL shows monotonic improvement), but conflates data volume with data diversity—adding synthetic data both increases sample count and adds domain coverage. A random-subset study at fixed domain distribution would isolate the effect of data volume. If accuracy continues to improve with data volume, it would justify a large-scale data collection effort (e.g., scraping more repositories, generating more synthetic data with different models). If accuracy plateaus at 70k, it would redirect research toward architectural innovations (larger context windows, structured prediction, retrieval-augmented translation).

Practical Applications and Downstream Use Cases

Batch transpilation of legacy CUDA codebases for hardware migration planning. Organizations with large, stable CUDA codebases (scientific computing labs, financial modeling firms, physics simulation groups) considering a migration to AMD hardware can use the CASS source translation models (90–95% accuracy, Table 3) as a first-pass transpilation tool. The 7B model outperforms HIPIFY by 7.5% on CASS-Bench (95% vs. 87.5%, Table 3), and while this is measured on Claude-generated code rather than real-world CUDA, it suggests the model handles edge cases—preserving string constants, kernel launch syntax, output stream selection (Appendix A.5.2)—that HIPIFY's rule-based approach mishandles. An organization could run the CASS-7B model over their entire CUDA codebase, compile the resulting HIP code on AMD hardware, and use compilation failures to identify the remaining CUDA features that need manual attention, substantially reducing the initial triage cost compared to running HIPIFY alone. The 1.5B model (90% source accuracy) may be a more practical choice for organizations without access to 4× A100 GPUs, since it can presumably run on more modest hardware and still substantially outperforms HIPIFY.

Targeted generation of aligned assembly pairs for domain-specific fine-tuning. The CASS pipeline (Section 3) is designed as a general method for generating aligned GPU assembly pairs given any compilable CUDA source. An organization with a specialized GPU workload (e.g., seismic processing kernels, computational fluid dynamics solvers, custom ML operators) could: (1) collect their existing CUDA kernels, (2) run them through the CASS pipeline (HIPIFY conversion → dual-stack compilation → assembly extraction) using the provided open-source code (GitHub repository), (3) fine-tune the CASS-7B model on their domain-specific aligned pairs (perhaps 500–2,000 additional samples), and (4) achieve substantially higher assembly translation accuracy on their specific workload than the 37.5% general average. This is feasible because the pipeline is automated and the fine-tuning recipe is documented (Section 5: LLaMA-Factory, DeepSpeed, 16K context, 1e-5 learning rate). The paper shows that domain matters enormously—0% on math/graph tasks, 100% on physics simulations (Section 6, Figure 5)—and a lab with physics simulation kernels could expect high accuracy even from the base CASS models.

GPU architecture exploration and co-design using translated assembly as a proxy. Hardware architects and compiler engineers designing future GPU ISAs or evaluating the portability of proposed instruction set extensions can use the CASS assembly translation model as a rough "compatibility probe." The experiment: take a set of CUDA kernels that exercise a proposed new Nvidia ISA feature (e.g., a new memory access pattern or synchronization primitive), compile them to SASS targeting a hypothetical future architecture (using NVIDIA's ptxas with appropriate flags), and run the CASS assembly translator to see whether the resulting RDNA3 assembly preserves the computational structure. If the model produces plausible RDNA3 that retains the memory access pattern of the SASS source, it suggests the feature has a natural AMD analog and that code using it would be portable. If the model produces degraded or nonsensical assembly, it suggests the feature is NVIDIA-specific and portability would require manual intervention. This is not a precise engineering tool—37.5% accuracy is far too low for that—but it provides a rapid, automated first-pass signal at a stage of the design process where manual translation of hundreds of kernels would be prohibitive. The execution fidelity findings (85% of correct translations within ±5.6% of native performance for runtime and memory, Appendix A.4.2) provide some confidence that when the model does produce correct assembly, it preserves performance characteristics, making the proxy somewhat informative.