# SCALE by Spectral Compute

> SCALE is a GPU compiler toolkit that natively compiles unmodified CUDA source code for AMD and NVIDIA GPUs. It is a drop-in replacement for nvcc with multi-vendor GPU support. SCALE enables running CUDA on AMD GPUs without code rewrites, eliminating NVIDIA vendor lock-in. Made by Spectral Compute.

SCALE compiles a superset of CUDA — including inline PTX asm — directly to native AMD GPU machine instructions and to NVIDIA PTX instructions. No emulation, no translation layer, no runtime overhead. It is an enterprise-level solution that delivers vendor optionality, performance and increased productivity in CUDA development via novel compiler optimizations and better-than-native developer tooling. It is a direct alternative to HIPIFY, SYCLomatic, ZLUDA, and porting to a different programming language.

This file collects the content of https://scale-lang.com in plain text. The short index is https://scale-lang.com/llms.txt; technical documentation is at https://docs.scale-lang.com/stable/.

## What is SCALE?

Decoupling Code from Silicon.

CPU developers don't rewrite their software for every new chip architecture—they simply recompile. SCALE brings this standard of portability to GPU computing.

It is a comprehensive toolkit—combining a cross-compiler, drop-in libraries, and language extensions—that acts as an agnostic interface between your code and the hardware.

With SCALE, you can take your existing HPC applications (starting with CUDA) and deploy them to your accelerated compute platform of choice.

- Your CUDA codebase is the single source of truth with zero porting required.
- Unlock new efficiency gains through advanced software optimization, not just hardware upgrades.
- Access enhanced UX and language extensions built by HPC developers, for HPC developers.
- Bring vendor neutrality to your CUDA code: choose hardware based on cost, performance or availability.

## Licensing

Free for non-commercial purposes. Paid license for commercial use; available for design partnership and support.

- **Research & Non-Commercial**: For non-commercial, educational, and research purposes on all client, workstation and data-center GPUs.
- **Commercial**: Standard license for commercial deployment and use. Contact us for pricing.
- **Enterprise**: Collaborate with our team on custom solutions, optimizations, dedicated support, and roadmap prioritization.

Contact: hello@spectralcompute.com

## Frequently asked questions

### Is SCALE free?

SCALE is free for non-commercial use including research and academia. For commercial use, a license agreement is required. Read more [here](https://scale-lang.com/#licensing).

### Does SCALE support PyTorch?

Yes. PyTorch works with SCALE and is available as a pre-release on request — [get in touch](https://scale-lang.com/solutions#contact) to get access. For a detailed overview of the currently supported CUDA projects, see [this table](https://github.com/spectral-compute/scale-validation?tab=readme-ov-file#current-status) of our validation suite.

### Does SCALE support my GPU?

SCALE supports a wide range of both consumer and enterprise GPUs, and will support more in the future. For a detailed overview, see [this section](https://docs.scale-lang.com/stable/#which-gpus-are-supported) of the official SCALE documentation.

### Does SCALE increase performance?

In many cases, yes, it does. Reducing compute costs can be a good reason to choose SCALE. For the latest performance benchmarks, see [this section](https://scale-lang.com/#benchmarks) of our website.

### <Vendor> has <Solution>, why should I use SCALE instead?

SCALE is centered around CUDA and allows you write your code once, and run everywhere with zero code rewrite. It is a drop-in replacement for `nvcc`. For full explanation of all the differentiators of SCALE, see [this section](https://docs.scale-lang.com/stable/manual/comparison/) of our technical documentation.

### Does SCALE infringe NVIDIA’s copyright?

By design, SCALE does not infringe NVIDIA’s EULAs or copyright. We think CUDA is amazing and we follow the guidelines set by NVIDIA. Check out [this post](https://www.linkedin.com/posts/spectral-compute_write-cuda-run-everywhere-activity-7399130853668245504-dvWQ?utm_source=share&utm_medium=member_desktop&rcm=ACoAAFgca6cBto8wrYuRnLPzR6DHSH67dicqjbE) for more information.

## Docs

- [SCALE Documentation](https://docs.scale-lang.com/stable/): Official documentation, installation guide, API reference, and examples
- [Install SCALE](https://docs.scale-lang.com/stable/manual/tutorials/how-to-install/): Installation instructions for Ubuntu, Rocky Linux, tarball, and Docker
- [Comparison to Other Solutions](https://docs.scale-lang.com/stable/manual/comparison/): Detailed comparison of SCALE vs HIP and ZLUDA
- [CUDA Dialects](https://docs.scale-lang.com/stable/manual/dialects/): Explanation of the nvcc vs Clang-CUDA dialect problem that SCALE solves
- [API Coverage Report](https://docs.scale-lang.com/stable/manual/apis/): Current CUDA API coverage status
- [Language Extensions](https://docs.scale-lang.com/stable/manual/language-extensions/): SCALE's opt-in extensions beyond standard CUDA
- [Inline PTX Support](https://docs.scale-lang.com/stable/manual/inline-ptx/): How SCALE compiles inline PTX asm for AMD GPUs
- [FAQ](https://docs.scale-lang.com/stable/manual/faq/): Frequently asked questions

## Examples and Validation

- [Validated CUDA Projects](https://github.com/spectral-compute/scale-validation): 53+ open-source CUDA projects tested with SCALE, including Llama.cpp, GROMACS, CUTLASS, Faiss, and XGBoost
- [Basic Example](https://docs.scale-lang.com/stable/examples/basic/): Simple CUDA program compiled with SCALE
- [PTX Example](https://docs.scale-lang.com/stable/examples/ptx/): Inline PTX asm example
- [BLAS Example](https://docs.scale-lang.com/stable/examples/blas/): cuBLAS usage with SCALE
- [Performance Improvement Example](https://gist.github.com/ChrisKitching/53d9b3503c11881c5ea395ec41aff419): Managed memory microbenchmark (HIP vs. SCALE)
- [Benchmarks](https://scale-lang.com/#benchmarks): Extended OpenDwarfs results: on par with HIP and nvcc on average, and faster on about half the workloads
- [Rodinia Benchmarks on MI355X](https://scale-lang.com/benchmarks/rodinia-mi355x-february26): SCALE vs HIP on an AMD Instinct MI355X, February 2026

## Company

- [Company Website](https://spectralcompute.com): The company behind SCALE
- [Solutions](https://scale-lang.com/solutions): SCALE for silicon vendors, cloud service providers, and HPC enterprises
- [LinkedIn](https://www.linkedin.com/company/spectral-compute): Official company page
- [Discord](https://scale-lang.com/s/discord?utm_source=llms_txt): Community and support
- [X (Twitter)](https://x.com/SpectralCom): Latest updates and news
- [r/CUDAUnlocked](https://www.reddit.com/r/CUDAUnlocked/): Subreddit for running CUDA on any GPU
- Sales: hello@spectralcompute.com
- General inquiries: hello@spectralcompute.com
- CEO: ms@spectralcompute.com

# Blog posts

Technical posts on CUDA portability, benchmarks, and developer experience from the team behind SCALE.

---

## Why hardware-agnostic isn't the same as lowest-common-denominator

URL: https://scale-lang.com/posts/2026-06-12-hardware-agnostic-vs-lcd
Published: 2026-06-12
Author: Michael Søndergaard

*Part 4 of a series on 'why Spectral exists', \~10 minute read.*

[Part 3](https://scale-lang.com/posts/2026-05-27-cuda-was-always-cross-platform) argued that the CUDA programming model is more general than the hardware it grew up on, and that the gap between the hardware it can target and the hardware it was designed for is the work. That argument has a cost attached to it, and a careful reader will already be reaching for the bill.

The objection goes like this. The moment you commit to running across hardware from more than one vendor, you have committed to expressing your program in terms of what those vendors have in common. Anything one architecture can do that the others can't is off the table, because the abstraction has to hold for all of them. So a cross-vendor system is a lowest-common-denominator system by construction, and lowest-common-denominator means slow. You can have portability or you can have peak performance. SCALE has chosen portability, and the performance tax is the price.

This is the most serious thing a technical investor or a competitor will say about what we do, and they are right about a whole category of systems. They are wrong about ours, and the reason they're wrong is the entire engineering bet. This post is about the difference.

### **Two things called portable**

The word "portable" covers two mechanisms that have almost nothing in common, and the objection is true of one of them.

The first is portability by abstraction. You define a layer that sits above the hardware, you target the layer, and at runtime something maps the layer onto whatever machine you happen to be on. An interpreter does this. A runtime shim does this. A naive source-to-source translator does this. The layer can only express what every target supports, because it has no knowledge of the specific machine at the point where the program is written — that knowledge arrives too late to change the code. This is genuinely lowest-common-denominator, and it genuinely costs performance. The objection describes it perfectly.

The second is portability by compilation. You transform the program into a representation that preserves its structure, and then you generate native code for each target separately, with full knowledge of that target, ahead of time. The representation is shared; the code that comes out the other end is not. Nothing about the program is pinned to the intersection of what the targets have in common, because the decisions that matter are made per target, after the target is known.

These two things get filed under the same word, and the objection quietly assumes the first when the system in question is the second. SCALE is the second. Everything else in this post follows from that.

### **The LLVM proof**

The cleanest way to see that cross-vendor and native-quality aren't in tension is to look at a place where the argument has already been settled, by twenty years of production use, in a system nobody accuses of being slow.

C compiled through LLVM runs on x86, ARM, and RISC-V. It is competitive with hand-written assembly on all of them. Nobody describes LLVM as a lowest-common-denominator layer, even though it is, by definition, an abstraction over multiple instruction set architectures — which is exactly the property the objection claims must cost performance. The LLVM IR does not express the intersection of what x86 and ARM can do. It expresses the structure of the computation, and then each backend generates native code with full knowledge of its own instruction set, register file, cost model, and scheduling constraints. The differences between the architectures simply shift into the realm of code generation.

That is the whole trick, and it is not a new one. Cross-architecture and native-quality coexist when the abstraction sits at the right level and the target-specific decisions are made by a backend that knows the target. SCALE is the same category of system, one layer up. The CUDA source is the high-level description; the per-vendor backend is where the machine gets exploited. If you accept that LLVM doesn't make C slow, you have already accepted the mechanism that lets a compiler be cross-vendor without being lowest-common-denominator. The only open question is whether we've built the backends well, which is an engineering question, not a question about whether the approach can work at all.

### **Where the objection is right**

It is worth being precise about when "abstraction costs performance" is true, because it often is, and pretending otherwise would be the kind of marketing claim this series exists to avoid.

The tax is real whenever the target is decided at runtime rather than compile time. A dynamic translation layer that rewrites instructions on the fly pays for the translation and can't afford the expensive analysis that good code generation needs. A thin runtime shim that forwards calls to whatever device is present has no opportunity to specialise, because it doesn't know what it's running on until it's already running. Anything that targets the common subset by construction — because it was designed as a portability layer first and a performance system second — inherits the ceiling the objection describes.

What these have in common is the moment the target becomes known. If the program is committed to its final form before the machine is known, the machine can't inform the code, and the common-denominator ceiling is real. SCALE makes its decisions at compile time, per target, with the target fully in hand. Fundamentally, this differs from a runtime shim for strictly structural reasons: it is all about when the vendor-specific knowledge gets to act on the code. For us, it acts before a single instruction is emitted.

### **What a compiler does that a translator can't**

Here is the claim the objection gets exactly backwards. Vendor-specific optimization is not erased by the cross-vendor model. The cross-vendor model is *where the vendor-specific optimization happens*. A translator substitutes one vendor's concepts for another's and stops. A compiler re-derives the right code for each machine from scratch, and the machine it's deriving for is the one it specialises against.

Concretely, when SCALE targets AMD, the backend is not looking up the AMD spelling of an NVIDIA concept. It is reasoning about AMD's hardware:

- **Occupancy is computed against the actual target.** The number of wavefronts a compute unit can keep in flight depends on register and LDS pressure, and the tradeoff curve is AMD's, not NVIDIA's. The backend allocates registers and schedules against AMD's curve, which is a different optimisation problem with a different answer, not a translated version of the NVIDIA one.  
- **Shared memory and LDS decisions are made for the target's sizes and banking.** What lives in fast on-chip memory versus what stays in registers is a target-specific call, and the banking behaviour that determines whether those accesses conflict is AMD's banking behaviour.  
- **The wavefront width is part of the cost model, not a find-and-replace.** A warp is 32 lanes and a wavefront is commonly 64\. That changes how reductions are shaped, how divergence costs, and how memory accesses coalesce. The backend knows the width and generates code that's right for it.  
- **Memory access patterns are laid out for the target's coalescing rules.** What counts as a coalesced access — and therefore how the compiler should arrange loads and stores — is a property of the specific memory subsystem, and the backend lays out for the one it's compiling for.

None of this is available to a translator, because a translator's unit of work is the concept, and these decisions live below the concept, in the code generation. The reason a compiler can match native code is that it is *doing* what the native programmer would do: making target-specific decisions with target-specific knowledge. The cross-vendor part is upstream of all of it.

### **The evidence**

This stops being an argument the moment there's a number, because the objection makes a falsifiable prediction. A lowest-common-denominator system cannot match a hand-targeted native port. That's what lowest-common-denominator means — it's leaving the vendor-specific performance on the table by construction. So if SCALE matches a native HIP port on a real workload, the system is not lowest-common-denominator. The parity is the proof.

It does match. On GROMACS, unmodified CUDA compiled through SCALE runs on AMD at performance on par with the native HIP port, and on the [MI355X work with TensorWave](https://tensorwave.com/blog/spectral-compute-runs-unmodified-cuda-workloads-on-tensorwaves-infrastructure-with-up-to-25-7x-faster-performance) the picture is the same — native-quality code out of source that was never written for the target. A common-denominator layer is incapable of reaching parity with code that was hand-targeted to the machine. Reaching it anyway is the empirical refutation of the objection, not a rhetorical one. The benchmark and the claim are the same fact stated two ways.

### **The harder version of the objection**

The sophisticated version of the pushback concedes the parity and moves the line. Fine — you match the HIP port. But you'll never beat a kernel that an expert hand-tuned for one specific piece of silicon, so there's still a gap, and the gap is the tax.

Two answers. The first is that hand-tuned-per-kernel is not the baseline that real codebases run at. Almost nobody is hand-tuning every kernel for every target; they're shipping the port they could afford to write, and the comparison that matters in practice is against that, not against an idealised expert who tuned everything. Against the realistic baseline, the gap is frequently zero or negative.

The second is that closing the remaining gap against the idealised baseline is precisely what compiler technology has always done, given time. Auto-vectorised SIMD started well behind hand-written assembly and is now routinely as good or better, and it's portable across instruction sets in a way hand-written assembly never is. The same arc applies here. As [Part 3](https://scale-lang.com/posts/2026-05-27-cuda-was-always-cross-platform) put it for the tensor-core case, the hard transforms are fundamentally a pile of permutations to be cancelled or fused — complex, but tractable, and the kind of thing a compiler closes over time rather than a wall it hits. The gap is the work. It is not a ceiling, and a ceiling is what the objection needs it to be.

### **Why this matters for the thesis**

If the objection were right, the series collapses. A substrate that taxes performance cannot be the standard for accelerated computing, because nobody standardises on the thing that makes their code slower — the entire history of this field is people abandoning portable-but-slow for fast-but-locked-in the instant the performance gap becomes load-bearing. Native-quality code generation isn't a nice property of SCALE on top of the argument in Parts 1 through 3. It's the precondition for any of that argument to hold.

Part 3 said the model is general. This post says the generality is free — that you don't pay for cross-vendor reach with performance, because the compiler makes the per-vendor decisions a native programmer would make, with the same information. The model being general and the performance being native are not two separate claims that happen to both be true. They're the same claim, and the compiler is the thing that makes both of them true at the same point in the pipeline.

### **Closing**

Hardware-agnostic describes where the code runs. It does not describe how well it runs, and the assumption that the two are the same is an intuition imported from a different kind of system. The intuition is correct about runtimes, interpreters, and translators, where the target shows up too late to shape the code. It is wrong about compilers, where the target shows up exactly in time.

The lowest-common-denominator objection is the best argument against what we do, and it is an argument about a system we didn't build. The model is general, the performance is native, and the compiler is where both of those stop being a tension and start being the same sentence.

Still not convinced? Join [our Discord](https://scale-lang.com/s/discord?utm_source=blog). We love talking compilers, GPU programming, and having exactly these sorts of debates.

---

## CUDA was always cross-platform

URL: https://scale-lang.com/posts/2026-05-27-cuda-was-always-cross-platform
Published: 2026-05-27
Author: Michael Søndergaard

*Part 3 of a series on ‘why Spectral exists’, \~10 minute read.*

[Part 2](https://scale-lang.com/posts/2026-05-19-the-brain-still-needs-the-hammer) argued that the compiler stack matters more in the agent era, not less, and that the substrate that agents write against should stay stable while the hardware varies. There's a hole in that argument I didn't address. It assumes CUDA can be the stable substrate even across hardware that looks fundamentally different from what it was designed for. A reasonable reader could grant everything in Parts [1](https://scale-lang.com/posts/2026-05-13-how-cuda-won-by-not-being-a-standard) and [2](https://scale-lang.com/posts/2026-05-19-the-brain-still-needs-the-hammer) and still ask: sure, but why CUDA? Isn't it just NVIDIA's hardware description language with extra syntax?

That's the question this post answers. The claim: the CUDA programming model is more general than the hardware it grew up on, and this is not coincidental. For the past two decades, NVIDIA hardware has changed significantly, yet the core parts of CUDA really haven’t.

### **The wrong mental model**

The mental model most people carry around is that CUDA is a thin layer over NVIDIA hardware. Threads map to CUDA cores, warps map to scheduler quanta, shared memory maps to a specific SRAM block on the streaming multiprocessor. The abstraction, in this view, is the hardware with a C++ syntax glued on top.

If that were true, porting CUDA to anything else would be a translation exercise. Find the AMD equivalent of each NVIDIA concept, substitute it in, and call it a day. The ceiling of the approach would be hardware that already looks like an NVIDIA GPU. Anything architecturally different — TPUs, systolic arrays, neuromorphic chips, whatever comes next — would be outside the model's reach by construction.

It isn't true. That's the wrong mental model, and almost everything interesting about CUDA's longevity follows from why it's wrong.

### **The right mental model**

CUDA is a way of describing parallel computation. Specifically: a hierarchy of parallel work (grid, block, thread), an explicit hierarchy of memory locality (global, shared, register), and an execution model — [SIMT](https://en.wikipedia.org/wiki/Single_instruction,_multiple_threads) — that lets the programmer reason about parallelism without micromanaging the hardware underneath.

None of those concepts are NVIDIA-specific. They're statements about how parallel computation decomposes. The grid-block-thread hierarchy is a way of expressing "this work is parallel across the whole problem, this work is parallel within a group that can share state, this work happens inside a single worker" — a structure that exists in any massively parallel computation, on any hardware that does massively parallel computation. The memory hierarchy is the same: every accelerator architecture has the property that some memory is close and fast and small, some is far and slow and large, and the programmer's life is easier if the model exposes that distinction explicitly rather than pretending it doesn't exist.

The hardware is the implementation of the model. The model is separable from it. This is the same move [LLVM IR](https://llvm.org/docs/LangRef.html) made twenty years ago — define the abstraction at the right level, and the backend becomes a code generation problem rather than a translation problem. The reason LLVM works across architectures isn't that x86 and ARM are secretly the same. It's that the IR was designed at a level of abstraction where the differences become a matter of code generation, not a matter of conceptual mismatch.

CUDA's longevity comes from the same property. The model was, by design, pitched at a level of abstraction general enough to outlive the specific GPU it was born on.

### **AMD is the easy case**

At the ISA level, the gap between AMD and NVIDIA GPUs is comparable to the gap between different generations of NVIDIA GPUs. CUDA was designed to abstract away the hardware differences within NVIDIA's own catalogue — and that machinery turns out to be exactly the machinery you need for cross-vendor portability:

* **Different [warp sizes](https://docs.scale-lang.com/stable/manual/compiler/diagnostic-reference/#-wptx-wave64)** are handled by the `warpSize` constant, and SCALE's [compiler diagnostics](https://docs.scale-lang.com/stable/manual/compiler/diagnostic-reference/#-wptx-wave64) to catch mistakes.  
* **Different shared memory sizes** are handled by occupancy APIs, launch bounds, and device property queries. NVIDIA already ships devices ranging from 64KB to 220KB, so CUDA programs deal with this already.  
* **Different occupancy characteristics** are handled by the same occupancy APIs already required to get optimal results across NVIDIA's own devices.  
* **Low-level microarchitectural details** sit below the programming model, where the compiler backend handles them.

None of this is new machinery added for AMD support. It's CUDA being used the way it was always meant to be used. Programs that aren't using these features will benefit from starting to, on NVIDIA hardware too.

Tensor cores are the exception. They aren't abstracted by the programming model at all, and they're radically different between vendors — and between NVIDIA generations. But even here, the big picture is the same: a warp or group of warps performs a matmul in registers. The layout of matrix elements across threads and the size of the matmul differ; the fundamental operation doesn't. This is the same situation CPU compilers faced when SIMD instruction sets started proliferating, and the resolution was the same one that has to happen here: build the compiler technology to model the operation and pattern-match naturally-written code onto it. Auto-vectorised SIMD CPU code is now routinely as good as or better than hand-written, and it's portable. The equivalent transform for GPU matmul accelerators is complex but tractable — fundamentally a pile of permutations to be cancelled out or fused into memory ops.

AMD is the proof of concept, not the proof of generality. If the model only ported to architecturally similar hardware, the skeptic would still be right that CUDA is a GPU-shaped abstraction. The interesting case is the one where the hardware doesn't look like a GPU at all.

### **The interesting case: GEMM on a systolic array**

The workload that pays for everything in modern AI is matrix multiplication. So start there.

A well-written CUDA GEMM kernel does something specific. It tiles the output matrix into blocks, assigns each block to a thread block, and has the thread block cooperatively load tiles of the two input matrices into shared memory. Threads within the block then do a register-level accumulation, multiplying their share of the loaded tiles into a partial sum that lives in registers. When the block finishes accumulating across the K dimension, the result tile is written back to global memory. This is the CUTLASS decomposition, and it's the basic shape of every serious GEMM kernel in the wild.

Now look at how a systolic array does the same thing. A systolic array — the design that defines TPUs and most contemporary AI accelerators — is a fixed grid of multiply-accumulate units, wired together so that operands stream through the grid in a coordinated pattern. The A matrix streams in from one edge, the B matrix from another, and each processing element receives operands from its neighbours, multiplies them, accumulates into a local register, and passes the operands along to the next PE. The output accumulates in place. The entire matmul happens as a pipelined dance across the grid.

These look like fundamentally different mental models. They aren't.

The CUDA block becomes the systolic tile. The shared memory load becomes the edge feed into the array. The thread-level multiply-accumulate becomes the PE-level multiply-accumulate. The K-dimension loop in the kernel becomes the pipeline depth of the systolic pass. The same logical decomposition — load a tile, compute on it, write the result — describes both the GPU kernel and the systolic array execution. The CUDA source is describing a structure of computation, and the systolic array is one possible *schedule* for that structure, where the GPU is another.

This isn't a clever observation about computation in the abstract. The mapping from a tiled matmul to a systolic schedule is well understood — XLA implicitly relies on it, and every TPU compiler in the world does some version of it. The catch is the direction. TPU compilers don't start from CUDA. They start from tile-level abstractions designed to lower cleanly onto systolic hardware — [Pallas](https://docs.jax.dev/en/latest/pallas/design/design.html), [Triton](https://github.com/triton-lang/triton), etc. — because nobody has built the compiler IP to *raise* a low-level representation like CUDA up to something a systolic backend can target. The lowering pathway exists; the raising pathway doesn't. That's why TPUs and other exotic accelerators have ended up with completely separate programming models, even though the structure of the computations they run is the same as what a CUDA kernel describes.

So the question isn't whether the CUDA model can express systolic computation. It can. The question is whether the compiler can recover enough structure from CUDA source to schedule it onto hardware that doesn't look like a GPU. That's a much harder problem than lowering from a high-level tile representation — and it's the problem that has to be solved if the substrate is going to stay stable while the silicon varies.

The novelty is the next step.

### **The general claim**

If GEMM ports because the CUDA model describes structure rather than hardware, the same is true for other kernels — though the work the compiler has to do varies with the target. On SIMT-like hardware, the source already describes what the silicon does, and the raising stack is light. On something like a TPU, the compiler has to lift the source much further before it can schedule it onto the systolic execution model. Stencils, reductions, fused custom ops, attention kernels — the structural information is there in the CUDA source. The question is how much raising it takes to recover that structure for a given target, not whether the structure exists.

The compiler is what realizes this. The CUDA source expresses the structure of the computation at a level of abstraction high enough that the backend has room to be clever about how to execute it. This is the same property that lets LLVM turn serial C into vectorized SIMD: the source describes the computation, the compiler picks the implementation that fits the target. CUDA, properly understood, is in the same category. The model is a description of parallel structure, and the backend's job is to find the implementation that fits the silicon.

The implication matters. CUDA is not a candidate for the standard substrate of GPU programming. It's a candidate for the standard substrate of *accelerated computing*, full stop. The space of hardware the model can target is much larger than the space of hardware it was designed for, and the gap between those two spaces is the work.

NVIDIA itself has been extending the substrate in exactly this direction. CUDA 13.1, released at the end of 2025, introduced cuTile — a tile-based programming model that sits alongside SIMT, lets the developer describe the algorithm at a higher level, and has the compiler and runtime handle the partitioning onto threads. The underlying representation, CUDA Tile IR, is MLIR-based, in the same compiler-infrastructure family that Pallas, Triton, and the TPU stacks use. CUDA programmers now have both registers: SIMT when the workload wants explicit thread-level control, tile abstractions when the workload is more naturally expressed at the tile level. Same substrate, more ways to write into it. The argument of this post is that the CUDA *model* was always more general than the hardware it grew up on; cuTile is the model continuing to grow into that generality, on NVIDIA's own roadmap. Having the cake and eating it is the right outcome for a substrate that intends to outlast any specific generation of silicon.

The obvious objection is that other candidate substrates exist. OpenCL was designed for portability and has middling vendor support to show for it. SYCL is gaining traction in HPC but the corpus is small and the tooling is recent. Triton sits one level above CUDA and is excellent for kernel authoring, but it's an addition to the ecosystem rather than a replacement for the layer below. Mojo is new and unproven. The argument for CUDA isn't that it's *uniquely* well-designed as an abstraction — several of the alternatives are arguably cleaner in some areas. But in addition to serving NVIDIA’s hardware evolution well for 20 years, CUDA has the corpus, the muscle memory of every working GPU programmer, and the largest training distribution for any agent that will generate kernels in the future. The substrate question is partly technical and mostly path-dependent.

### **Why this matters for the thesis**

Parts [1](https://scale-lang.com/posts/2026-05-13-how-cuda-won-by-not-being-a-standard) and [2](https://scale-lang.com/posts/2026-05-19-the-brain-still-needs-the-hammer) made claims about *what* the substrate should be and *who* should provide it. This part makes the claim about *why* CUDA specifically deserves that role. The model is general, the corpus is large, the semantics are stable. If you were going to pick a programming model to be the standard across all accelerated computing — across GPUs from every vendor, across systolic arrays, across whatever the next generation of AI silicon turns out to look like — you would want one with exactly those properties. CUDA already has all three. No other candidate has them at the same scale.

The compiler is where the generality gets realized. Without a compiler that knows how to target each architecture, the model's generality is a theoretical property — interesting to think about, useless in production. With one, it's the basis for a future where the code people have already written, the kernels agents have already learned to generate, and the muscle memory accumulated across two decades of GPU programming all keep working as the hardware underneath changes shape.

### **Closing**

CUDA is not a hardware description language. It's a programming model that happened to be born on one piece of hardware and has spent two decades being refined into something far more general than its origins suggest. The shape of AI hardware over the next decade is going to surprise people. New vendors, new architectures, new ideas about how silicon should be organized to run the workloads that matter. Some of those bets will pay off and some won't. We don't know which.

What we do know is that whichever hardware wins, the people building on top of it will want their existing code to keep working, their agents to keep generating code they can trust, and their substrate to stay stable across the change. That's the future [SCALE](https://scale-lang.com/) is built to support — whichever vendor, whichever shape, whatever comes next. The model is general enough to get there. The work is making the compiler match.

---

## SCALE 1.7 is out

URL: https://scale-lang.com/posts/2026-05-20-scale-1-7-release
Published: 2026-05-20
Author: Giulio Malitesta

This month's SCALE release is a meaty one. It pushes on three fronts at once: fleet deployability for cloud providers, raw HPC performance, and the PyTorch coverage that AI shops need to take AMD seriously.

Head over to our [install page](https://docs.scale-lang.com/stable/manual/tutorials/how-to-install/?utm_source=blog&utm_medium=internal&utm_campaign=blog_scale_release_announcement) to get the new release, and let us know what you find. Bug reports and feedback are welcome on [our Discord](https://scale-lang.com/s/discord?utm_source=blog).

## Multi-architecture AMD builds

SCALE can now produce a single binary that targets multiple AMD GPU architectures at once, with the right code path picked at runtime - the same multi-ISA bundling HIP supports. One artifact, one validation pass, the whole AMD fleet.

This matters for anyone shipping a CUDA workload across a mixed AMD fleet:
+ For **HPC enterprises and Independent Software Vendors**, it collapses the per-SKU build matrix - one binary covers all their AMD targets.
+ For **neoclouds and inference-as-a-service providers** building their own runtime images, it does the same thing *plus* gives the scheduler full placement freedom across the fleet.
+ For **CSPs renting bare GPU-hours**, it makes their AMD catalogue materially easier for customers to actually adopt - especially the long-tail SKUs.

One last thing to note: SCALE provides first-class support for older AMD architectures no longer supported by ROCm. See [our documentation](https://docs.scale-lang.com/stable/#which-gpus-are-supported?utm_source=blog&utm_medium=internal&utm_campaign=blog_scale_release_announcement) for details. If you have any of that hardware in your fleet, SCALE keeps it earning.

## Kokkos support
As you may know, [Kokkos](https://github.com/kokkos/) is the C++ performance-portability library behind a huge fraction of DOE simulation codes, hosted by the Linux Foundation and developed primarily by [Sandia National Laboratories (SNL)](https://www.sandia.gov/), [Oak Ridge National Laboratory (ORNL)](https://www.ornl.gov/), and the [French Alternative Energies and Atomic Energy Commission (CEA)](https://www.cea.fr/english/Pages/Welcome.aspx). You write your computation once against Kokkos's abstractions, then pick a backend at compile time - CUDA for NVIDIA, HIP for AMD, and so on - and the same source runs on whichever hardware you targeted.

With SCALE 1.7, Kokkos's CUDA backend runs on AMD GPUs directly - no porting to the HIP backend, no second validation cycle. You keep the backend your code was written and validated for, and SCALE handles the AMD targeting underneath. Note that this is the initial support release and we are still ironing out bumps in collaboration with the Kokkos maintainers.

For national labs and HPC enterprises with existing Kokkos+CUDA codebases, this removes the largest single barrier to evaluating AMD hardware on real workloads.

## Managed memory: over 100x faster on HPC workloads

Substantial improvements to managed memory put SCALE more than 100x faster on many HPC workloads - up to an order of magnitude ahead of HIP on those same applications. Our CTO Chris has [published a small self-contained microbenchmark](https://gist.github.com/ChrisKitching/53d9b3503c11881c5ea395ec41aff419) - a single CUDA file that builds under both HIP and SCALE - that isolates the cost of managed-memory page migration so you can reproduce the difference yourself. On AMD hardware SCALE relocates those pages dramatically faster than HIP: on an MI355, the first iteration drops from roughly 648,000 ms under HIP to about 6,200 ms under SCALE (**yes, you read correctly: 10.8 minutes vs 6.2 seconds**). It's a deliberately simple test, and even on a purely unidirectional transfer the gap is stark; workloads that alternate between host and GPU use would feel it even more. Our documentation covers Managed Memory more in detail [here](https://docs.scale-lang.com/stable/manual/library/performance-tuning/#managed-memory?utm_source=blog&utm_medium=internal&utm_campaign=blog_scale_release_announcement).

If you're running scientific code that lean heavily on unified/managed memory - CFD, molecular dynamics, climate, reservoir simulation - this feature alone makes this release worth picking up. We'll be publishing fresh benchmarks soon.

## What PyTorch needs

Three big steps toward making PyTorch run cleanly on SCALE:

- **Full cuBLASLt coverage.** The fused-GEMM hot path for modern transformer inference and training: GEMM+bias+GELU, and the rest of the fused-epilogue surface PyTorch leans on. On transformer workloads, this is where most of your GPU-hours get spent - a surface where hipBLASLt coverage is still maturing.
- **All the Graph API that PyTorch needs.** CUDA Graphs are how PyTorch captures and replays sequences of GPU operations with low launch overhead - important for inference and many training workloads.
- **Most of the cuSOLVER that PyTorch needs.** cuSOLVER is NVIDIA's dense linear algebra library (LU, QR, Cholesky, eigensolvers), and PyTorch leans on it for a fair amount of its linear algebra surface.

We're not done with PyTorch yet, but we're a meaningful chunk closer.

## OpenGL interop

SCALE now supports CUDA's OpenGL interop APIs - `cudaGraphicsGLRegisterBuffer`, `cudaGraphicsGLRegisterImage`, and friends. These let CUDA kernels and OpenGL share GPU resources without round-tripping through CPU memory.

This unblocks a category of CUDA projects that combine compute and rendering in the same pipeline: scientific visualization, particle and fluid simulations, real-time graphics tools, and a long tail of HPC/visualization software that previously had no AMD path at all.

It's also foundational groundwork for upcoming TMA (Tensor Memory Accelerator) support - the texture-descriptor and multi-dimensional addressing infrastructure that OpenGL interop relies on overlaps significantly with what TMA needs.

## FFmpeg on AMD
[FFmpeg](https://ffmpeg.org/) now builds and runs on SCALE, which means its CUDA-accelerated filters work on AMD GPUs. Media transcoding, scaling, and filtering pipelines built around NVIDIA's CUDA filters now have a drop-in AMD path - relevant for CDNs, video platforms, and any GPU-accelerated media workload looking to diversify off a single vendor.

## Other compiler optimisations

Several new compiler optimisations across the board, applying to both NVIDIA and AMD targets. Worth singling out: **shuffle and reduction improvements on NVIDIA targets, including automatic use of the `redux` instruction** - something NVIDIA's own documentation tells programmers to apply by hand. SCALE does it for you on the recompile.

Individually none of the rest is dramatic, but on long-running HPC and inference workloads, the compounding is the point.

## Bug fixes

The usual mix of fixes, including one worth calling out: [llm.c](https://github.com/karpathy/llm.c) now works on SCALE. You can see it running [in this video](https://youtu.be/Mg2miUydZQk?si=n7ee0jds6t8DF2Sk). [Alien](https://github.com/chrxh/alien) support has also been fixed, and now the CUDA-powered artificial life simulation program works across all supported AMD GPUs via SCALE.

---

## The brain still needs the hammer: Why compilers matter MORE in the agent era, not less

URL: https://scale-lang.com/posts/2026-05-19-the-brain-still-needs-the-hammer
Published: 2026-05-19
Author: Michael Søndergaard

_Part 2 of a series on why Spectral and SCALE exists._

In [Part 1](https://scale-lang.com/posts/2026-05-13-how-cuda-won-by-not-being-a-standard), I argued that cross-vendor portability in accelerated computing must be delivered by a company, rather than a committee, because the implementation is the standard. A reasonable reader would finish that argument by asking the obvious follow-up: fine, but what happens when the implementation gets written by an LLM?

It's a timely question. If agents can write CUDA, the skeptic says, then who needs a CUDA toolchain that runs everywhere? Just point the model at each backend and have it emit native code. CUDA here, ROCm there, SYCL there, etc. The compiler becomes a quaint historical artifact — the assembly programmer of the 2030s.

Let us take that view seriously, because at first hearing it sounds like it dissolves the whole thesis behind Spectral. It doesn't. The opposite, in fact: agentic code generation makes the compiler and runtime stack more valuable, not less. Here's why.

# LLMs and compilers do categorically different jobs

A compiler is a deterministic function from source to machine code. Same input, same output, every time, forever. That property is not incidental. It's the entire reason you can trust a billion lines of code running on hardware you've never seen.

An LLM is a probabilistic function from context to plausible tokens. Same input, different output, depending on temperature, sampling, model version, and the phase of the moon. That property is also not incidental — it's what makes the model useful for ideation, exploration, and synthesis. But it is exactly the wrong property at the metal layer.

Nobody wants merely stochastic correctness on FMAs, memory fences, or atomics. Nobody wants a kernel which produces subtly different results across runs because the model got creative about which warp did the reduction. Think of the LLM as the brain and the compiler as the hammer: the brain decides what to compute, the hammer forges the runtime executable, the same way, every time.

These two things look superficially similar — both produce code — but they sit on opposite sides of a categorical divide. Conflating them is the error.

# What agents actually need from a substrate

Once you accept that the agent and the compiler do different jobs, a useful question becomes: what does the agent need from the substrate — the compiler, runtime, and semantics it writes against — in order to be productive?

Three things, mostly.

*First*, fast and structured feedback. An agent iterating on a kernel needs compile errors it can parse, deterministic failure modes, and reproducible builds. The faster and more legible the feedback loop, the fewer attempts the agent will burn through before it lands on something which works. A toolchain emitting a wall of cryptic template errors is, for an agent, roughly as useful as no toolchain at all.

*Second*, one mental model rather than N (one for each vendor/GPU target). Every fragmentation of the substrate fragments the training distribution underneath the agent and multiplies the surface area for hallucination. If the agent has to remember that intrinsics are named one thing on NVIDIA, another on AMD, a third on Intel, and that the memory model has subtly different guarantees on each — its effective competence drops on every one of them. The model that deeply knows one substrate beats the model that knows N substrates shallowly.

*Third*, an environment that doesn't lie. The agent's productivity is bounded by how quickly it can iterate against something whose behaviour matches its documentation. A toolchain with silent miscompiles, undocumented edge cases, or platform-specific behaviour only showing up at runtime is a productivity sink whether the developer is human or not. Arguably more so when the developer is an agent, because the agent has no intuition for "that smells wrong, let me check."

So, not so surprisingly, these three things are also pre-requisites for human developer productivity.

# Volume changes the calculus

Here's the part that I think gets underweighted in conversations about agentic coding.

Human-authored GPU code is, in the grand scheme of things, a small corpus: tens of thousands of serious kernels, written mostly by specialists, reviewed mostly by other specialists. The substrate underneath those kernels can afford to be sharp-edged, because the people using it know where the edges are.

Agent-authored GPU code presents a fundamentally different volume of output. Potentially millions of kernels, increasingly directed by people who aren't GPU specialists and shouldn't have to be. The person building a vertical AI product is thinking about what their inference path should do. Whether memory accesses coalesce is precisely the kind of detail the stack should handle for them.

In that world, the substrate isn't a convenience. It's the load-bearing element. Without a robust compiler and runtime underneath the agent, you don't have 10x engineering output — you 10x the number of failed attempts it takes to get anything working. The thing that turns "generated code" into "running code" at scale is the toolchain. Take it away and the agent's apparent productivity collapses into a pile of broken kernels nobody can fix, generated by an agent that burned through a small fortune in inference tokens trying to fix them itself.

# CUDA is the substrate agents already know

Of all the GPU programming substrates in the world, CUDA has by a wide margin the largest public corpus of kernels, documentation, error messages, blog posts, and an accumulated archaeology of two decades of public internet Q&A. It also has the most stable semantics over the longest window. When you ask a frontier model to write a GPU kernel, it will be best at CUDA, by a considerable margin.

The "just emit native code for each backend" plan throws this away. It asks the model to do its worst-performing thing on every backend other than NVIDIA's, and it fragments the training distribution it does have. The agent ends up worse at all of them than it would be at the one.

The alternative is letting the agent write the substrate it actually knows well and let the compiler handle the targeting. The agent writes CUDA. SCALE makes it run on whatever's underneath. The agent's competence stays concentrated. The silicon stays interchangeable.

This isn't an argument about loyalty to NVIDIA's ISA, rather it's an argument about training corpora and semantic stability. If the world settled on SYCL tomorrow and had the models been trained on a decade of SYCL kernels, the argument would point at SYCL. It does not, because they were not.

# The compiler and runtime workload is growing

Step back from the agent question for a moment to look at the compiler and runtime stack as a workload. It is not shrinking. What people want from it now includes, but is not limited to:

 - Structured diagnostics designed to be consumed by an LLM rather than read by a human — schema-stable, machine-parseable, dense with the information an agent needs to fix the bug.
 - Incremental builds and tight iteration loops that make exploration cheap.
 - Autotuning surfaces the agent can drive end-to-end.
 - Profile-guided feedback the agent can act on.
 - Runtime introspection that exposes what's actually happening on the device in a form the agent can reason about.

Every one of these is a compiler-and-runtime problem. None of them get smaller as agents get better. They get larger, because the agent is now a first-class consumer of everything the toolchain emits, and the toolchain has to be designed accordingly.

The "compilers are obsolete" framing misses this part of the picture entirely. The agent does not replace the compiler. Instead the agent leans harder on it.

# Closing

Agentic AI doesn't make the compiler obsolete. It raises its leverage. The companies to win the agent era will be the ones whose substrate agents can target without ever thinking about the hardware underneath. Stable semantics, structured diagnostics, deterministic behaviour, one mental model that works everywhere — these are the properties that make agents productive, and every one of them is a compiler property.

The brain decides what to compute. The hammer forges the code to make it run.

With SCALE we are building the hammer and anvil to help you shape the intelligence of the future.

[Point your agent at it](https://docs.scale-lang.com/stable/manual/tutorials/how-to-install/) today and try it out.

---

## How CUDA Won by NOT being a Standard

URL: https://scale-lang.com/posts/2026-05-13-how-cuda-won-by-not-being-a-standard
Published: 2026-05-13
Author: Michael Søndergaard

_Part 1 of a series on why Spectral and SCALE exists, see [part 2](https://scale-lang.com/posts/2026-05-19-the-brain-still-needs-the-hammer)._

The way we standardize accelerated computing is broken, and it has been broken for long enough that the workaround: single-vendor capture, is now the default. It is worth saying out loud why the consortium model failed, what replaced it inside NVIDIA's walls, and what a cross-vendor version of the replacement actually looks like.

# Khronos was right about the problem

For thirty years, cross-vendor portability in graphics and parallel compute was delivered by consortia. A group of vendors would convene, draft a specification, ratify it, and then each member would ship a conforming implementation. OpenGL, Vulkan, OpenCL, SYCL — all variations on the same model. The Khronos Group is the canonical example, and to be clear: the diagnosis is right. If you want code to run across competing hardware, you need a neutral cross-vendor body. That part is not in dispute.

The mechanism, however, is a product of a different era.

The Khronos model works when the spec-ratification clock runs faster than the workload-evolution clock. Graphics APIs are a clean example. Triangles have not changed much in twenty years, so a spec ratified in 2016 still gets you most of what you need in 2026. The standardization clock ran ahead of the application clock, so the model held.

In accelerated computing, the clocks have inverted.

# Design by committee on an AI clock

The relevant workload for accelerated computing today is not graphics. It is dense linear algebra wrapped in increasingly exotic scheduling primitives, evolving on a release cycle measured in months. Flash Attention v1 to v3 took roughly two years and introduced fundamentally new programming patterns — warp specialization, asynchronous tile movement, persistent kernels. None of these patterns existed when SYCL 2020 was being drafted. They are unlikely to fit cleanly into SYCL 2030 either, because by then the patterns will have moved again.

A standards body cannot move at the pace of AI research. The structural reason is that a specification is, by construction, the intersection of what every vendor will agree to ship. Intersections move slowly because every member can veto. Meanwhile the workload moves at the pace of the fastest research team on Earth. In practice the mismatch is closer to 100x.

So in the field where standardization matters most, the historical mechanism is no longer fit for purpose.

# CUDA already proved the alternative

CUDA is not a standard. It has no consortium, no ratified spec, and no member companies. It is whatever NVIDIA ships in the next release of the toolkit. The "spec" is the implementation. The reference is the binary.

This is supposed to be a weakness — vendor lock-in, a single point of failure. It is all of those things. But it is also why CUDA wins. The CUDA team can introduce a new memory model, a new tensor core instruction, a new asynchronous copy primitive, and it ships. There is no committee meeting. There is no compatibility working group. The implementation moves at the pace of NVIDIA's hardware roadmap, which happens to match the pace of the workload, because NVIDIA is also the workload's largest direct customer.

Implementation-defined standardization is faster than specification-defined standardization. CUDA is the proof.

The catch is the obvious one: CUDA is one vendor's implementation. The standard is fast because it is captured. That is not an acceptable end state for the rest of the industry, and the rest of the industry knows it.

# The for-profit Khronos analogue

Spectral exists to resolve this contradiction. We preserve the right diagnosis from Khronos — that accelerated computing needs a neutral cross-vendor body — and replace the mechanism. Instead of a consortium that publishes a specification and waits for member implementations, we are a focused, funded engineering team that ships the implementation itself, for every vendor that wants to participate.

SCALE is that implementation. It compiles unmodified CUDA source — the de-facto programming model of the field — and produces native code for AMD and NVIDIA GPUs today, with additional architectures in flight. The "standard" is what SCALE compiles. The conformance test is whether your code runs, and runs fast. There is no document to argue about. There is a toolchain.

This makes Spectral structurally different from both a standards body and a chip vendor.

We do not build silicon, so we have no incentive to favour any architecture, and we do not run a consortium, so we are not bottlenecked on member agreement. We work directly with chip vendors, take feedback from the developers writing kernels in production, and ship the implementation that closes the gap on a timescale of weeks.

The for-profit structure is load-bearing here. A volunteer consortium can publish a PDF. It cannot maintain a production-grade compiler across multiple architectures at the pace AI demands. That requires a funded engineering team with hard deadlines and direct accountability to paying customers. Khronos identified the right role — the neutral cross-vendor body — and chose the wrong form. The form has to be a company.

# Why agents make this urgent

There is one more reason the specification model is finished, and the full force of it has not landed yet.

Code is increasingly written by language models. An agent that produces a CUDA kernel produces it in seconds. If you want that kernel to run on a non-NVIDIA target, the path from agent-generated source to executing binary has to operate at the same wall-clock speed. A spec-based portability story — wait for the consortium to define an equivalent primitive, then for the vendor to ship it, then for a developer to port the code — does not work when the producer of the code is operating at machine pace.

Implementation-defined portability does work. If the kernel compiles, it runs. If a new primitive shows up in the wild, the team shipping the implementation adds it. There is no third actor in the loop.

The pace of agentic code generation is going to make committee-driven standardization look the way ISO 9001 looks to a software startup: a process from a different era.

The question is not whether implementation-defined standards win. They have already won, inside one vendor.

The question is who builds the cross-vendor version.

We are. SCALE compiles unmodified CUDA code on AMD and NVIDIA today, with more architectures in flight: [Install SCALE now and try it out](https://scale-lang.com/)

---

## CUDA, the De Facto Standard of HPC?

URL: https://scale-lang.com/posts/2026-03-10-cuda-the-de-facto-standard-of-hpc
Published: 2026-03-10
Author: Giulio Malitesta

For nearly two decades, NVIDIA’s proprietary programming model, CUDA, has been the "gold standard" for unlocking the capabilities of the GPU. In June 2024, Jensen Huang announced that the ecosystem has reached 5 million developers[^1], growing ~150% from a 2M-developer userbase in 2020[^2]. This figure is a conservative floor; since registration isn't required on the NVIDIA download page and many install CUDA via package managers or Docker, the actual user base is likely much larger.

With [SCALE](https://scale-lang.com), **we treat CUDA as the de facto standard for HPC workloads**, allowing developers to compile their unmodified CUDA code for non-NVIDIA architectures. But this assumption is increasingly being challenged. New languages and paradigms - from **HIP** and **SYCL** to **Pythonic HPC approaches** - promise a future where developers are no longer locked into NVIDIA's hardware.

*Is this the beginning of a multi-vendor revolution, or is CUDA's lead unbeatable*? To find out, we did some research to probe the current developer landscape.

## The Contenders

| **CUDA (2006)**, developed by NVIDIA as the GPGPU first-mover. | **OpenCL (2009)**, proposed by Apple and maintained by the Khronos Group. | **SYCL (2014)**, a Khronos Group standard for modern, "single-source" C++. | **HIP / ROCm (2016)**, launched by AMD to challenge NVIDIA’s dominance. |
| :--- | :--- | :--- | :--- |
| Approachable "C++-like" syntax and a vast [library ecosystem](https://developer.nvidia.com/cuda/cuda-x-libraries) (cuBLAS, cuFFT). | Open, royalty-free standard for cross-vendor portability (NVIDIA, AMD, Intel). | Targets heterogeneous computing with host and device code in a single file. | Syntactic "dizygotic twin" to CUDA, allowing a shared codebase for NVIDIA and AMD. |
| Low barrier to entry: accelerate standard C89 code via libraries before moving to custom kernels. | Requires significant boilerplate code due to its low-level nature. | Unique design enables potential for novel compiler optimizations across the host/device boundary. | Includes [HIPIFY](https://github.com/ROCm/HIPIFY) to automate the porting of CUDA code directly to HIP, with limitations. |
| Proprietary and exclusive to NVIDIA hardware. | Often less attractive to developers compared to higher-level abstractions. | Core of Intel’s [oneAPI initiative](https://www.intel.com/content/www/us/en/developer/tools/oneapi/overview.html) for multi-vendor architecture development. | Positioned as a bridge for existing CUDA investments. |

## Methodology

### Raw Code Footprint via GitHub Search (Web UI)
First, we analyzed the volume of active codebases on GitHub using targeted search queries.
* **CUDA, OpenCL, and HIP**: We leveraged GitHub's native language classification and searched for [`language:CUDA`](https://github.com/search?q=language%3ACUDA), [`language:OpenCL`](https://github.com/search?q=language%3AOpenCL) and [`language:HIP`](https://github.com/search?q=language%3AHIP). It is important to acknowledge the concern that GitHub may incorrectly classify HIP code as CUDA; we address this overlap in later steps.
* **SYCL**: Because SYCL often resides in standard C++ files, we quantified its presence by summing results for core header inclusions, specifically [`#include <CL/sycl.hpp>`](https://github.com/search?q=%23include+%3CCL%2Fsycl.hpp%3E) and [`#include <sycl/sycl.hpp>`](https://github.com/search?q=%23include+%3Csycl%2Fsycl.hpp%3E). This approach may positively skew results as both core headers could appear in the same codebase.

### Intentional Tagging via GitHub Topics
To measure intentional developer alignment, we analyzed GitHub Topics. We counted repositories where developers explicitly tagged their project with a specific framework.

+ The specific queries used were: [`topic:cuda`](https://github.com/search?q=topic%3ACUDA), [`topic:opencl`](https://github.com/search?q=topic%3AOpenCL), [`topic:sycl`](https://github.com/search?q=topic%3ASYCL), and [`topic:hip`](https://github.com/search?q=topic%3AHIP).

### Idiom Analysis via GitHub REST APIs
Next, we searched for toolkit-specific idioms in active codebases. We used [`extract_cuda_without_hip.py`](https://github.com/spectral-compute/research-assets/blob/master/Blogpost%20-%20CUDA%2C%20the%20De%20Facto%20Standard%20of%20HPC%3F/extract_cuda_without_hip.py) and [`measure_alt_hpc_adoption.py`](https://github.com/spectral-compute/research-assets/blob/master/Blogpost%20-%20CUDA%2C%20the%20De%20Facto%20Standard%20of%20HPC%3F/measure_alt_hpc_adoption.py) to query for the syntax used to define or launch kernels. 

* **Code Idioms**: `cudaMalloc` (CUDA), `clCreateKernel` (OpenCL), `sycl::queue` (SYCL), and `hipLaunchKernelGGL` (HIP).

Because HIP codebases often contain standard CUDA API calls like `cudaMalloc` for cross-platform support, relying solely on keyword searches overcounts true CUDA lock-in in this specific case. To refine our research, `extract_cuda_without_hip.py` subtracts any repository containing AMD's `hip_runtime.h` from our baseline CUDA results. During data collection, occasional warnings indicating `"Bin [X] returned 1000+ files"` occur when a specific file-size interval hits GitHub’s hard limit of 1,000 results per query. However, because our methodology tracks unique repositories across over 150 different size bins, a project is only missed if every single one of its matching files falls into truncated bins. Consequently, these warnings indicate that the final repository count is a lower bound. The logs and output of each script can be found [here](https://github.com/spectral-compute/research-assets/tree/master/Blogpost%20-%20CUDA%2C%20the%20De%20Facto%20Standard%20of%20HPC%3F).

### Academic Mindshare via Semantic Scholar API
We developed a Python tool ([`academic_research.py`](https://github.com/spectral-compute/research-assets/blob/master/Blogpost%20-%20CUDA%2C%20the%20De%20Facto%20Standard%20of%20HPC%3F/academic_research.py)) to interface with the Semantic Scholar Graph API. We queried paper metadata for the years 2020 through 2025 to measure research momentum, signaled by any of the following strings:
* **For CUDA:** `"CUDA" | "NVIDIA GPU" | "cuBLAS"`
* **For HIP:** `"HIP" | "ROCm" | "AMD GPU" | "MI250"`
* **For OpenCL:** `"OpenCL" | "SPIR-V"`
* **For SYCL:** `"SYCL" | "OneAPI" | "DPC++" | "Intel GPU"`

This research is limited by the fact that generic terms like `"AMD GPU"` might capture papers that aren't specifically about the HIP programming language.

### Research scope and limitations
This research is limited to **publicly available codebases on GitHub** and **published academic papers**. We cannot investigate proprietary, "closed-source" enterprise codebases. Empirical business experience suggests that the dominance of CUDA in the private sector is even more pronounced; alternatives are often relegated to side-projects aimed at proof-of-concept (PoC) portability.

## Quantifying the Dominance

### GitHub: The "Developer Desert" for Alternatives
When we analyzed the raw code footprint on GitHub, the disparity was staggering. Below is the breakdown of the raw ecosystem metrics collected via the **GitHub Search UI**.

| Language/Framework | Files | File Share (%) | Repos | User Share (%) | Method |
|---|---|---|---|---|---|
| **CUDA** | 692,000 | 82.79% | 35,400 | 99.76% | Research query: `language:CUDA` |
| **OpenCL** | 92,700 | 11.09% | 0[^3] | 0.00% | Research query: `language:OpenCL` |
| **SYCL** | 37,600 | 4.50% | 0[^4] | 0.00% | Research query: sum of results for `#include <CL/sycl.hpp>` and `#include <sycl/sycl.hpp>` queries. |
| **HIP** | 13,600 | 1.63% | 85 | 0.24% | Research query: `language:hip` |

![Pie chart indicating 82.8% of GitHub Search metrics belong to CUDA](/assets/github_search_web_ui_chart.png)

We collected similar results by interrogating **GitHub Topics**. This metric represents repositories where developers explicitly tagged their project with a specific framework, showing intentional alignment. Interestingly, HIP retains the same percentage here as in the previous raw file search. This consistency suggests that GitHub’s potential misclassification of HIP code as CUDA may be less significant than generally assumed.

| Topic Tag | Repositories | Share (%) | Method |
|---|---|---|---|
| **CUDA** | 6,657 | 80% | "CUDA topic": `topic:cuda` |
| **OpenCL** | 1,326 | 16% | "OpenCL topic": `topic:opencl` |
| **SYCL** | 196 | 2.4% | "SYCL topic": `topic:sycl` |
| **HIP** | 134 | 1.6% | "HIP topic": `topic:hip` |

![Pie chart demonstrating 80.1% of GitHub Topics belong to CUDA](/assets/github_topics_chart.png)

Similar findings were observed by searching for indications of active codebases that sport specific **code idioms**.

| Signature | Language | Repos | Share (%) |
|---|---|---|---|
| `cudaMalloc` | CUDA | 59,950 | 77.1% |
| `clCreateKernel` | OpenCL | 11,887 | 15.3% |
| `sycl::queue` | SYCL | 3,355 | 4.3% |
| `hipLaunchKernelGGL` | HIP | 2,599 | 3.3% |

![Pie chart illustrating 77.1% repository share for the cudaMalloc code idiom](/assets/github_rest_api_idioms_chart.png)

### Academic Research
While researchers are increasingly interested in cross-vendor portability, results show the baseline volume of academic publications remains heavily skewed toward NVIDIA's ecosystem. Our analysis of paper count from 2020 to 2025 shows that **CUDA maintains an overwhelming 82.83% incidence** in academic research related to GPU programming and accelerated libraries. Academic interest in CUDA peaked at **193 papers in 2025 alone**, exceeding the sum of all alternatives considered.

| Metric | CUDA | HIP | OpenCL | SYCL |
|---|---|---|---|---|
| **Total** | 912 | 12 | 111 | 66 |
| **%** | 82.83% | 1.09% | 10.08% | 5.99% |

![Line chart representing the Semantic Scholar graph data, displaying a heavy volume lean towards CUDA](/assets/semantic_scholar_graph_api_chart.png)

While we continue to see case studies and benchmarks testing SYCL and HIP, the data suggests these represent a residual interest or a niche academic focus on portability rather than a ‘shifting tide’. The amount of academic interest for all alternative programming languages appears stable or even decreasing, while CUDA's dominance in research output continues to accelerate. However, because studies on the limits of performance portability offer valuable insights, we have compiled a few notable examples in the *Appendix* at the end of this post.

## The Economic Moat
Beyond the technical architecture, there is a powerful economic engine sustaining CUDA’s dominance. Our research suggests that the "moat" is built on three pillars of enterprise risk management:

* **The Talent Tax:** The most significant factor is the global workforce. With more than 5 million developers already proficient in CUDA, *the industry cannot afford the "talent tax" of retraining this massive population on new languages*. For most enterprises, accepting CUDA as the standard is a more sensible economic decision than funding a multi-year retraining program. Expertise is the hardest layer of the stack to swap.
* **Production Reliability and Sunk Cost:** Transitioning from CUDA to a different hardware abstraction effectively asks enterprises to *throw away their existing investments in validation and QA*. When a codebase is "production-proven" on CUDA, the perceived risk of moving to an unproven ecosystem - where subtle compiler bugs or performance regressions could derail a mission-critical project - is often too high to justify.
* **Saturation and the Failure of Alternatives:** With NVIDIA holding ~90%[^5] of the HPC infrastructure market, CUDA has effectively become the "source code" for expressing accelerated computing. *Many attempts to break this monopoly have inadvertently reinforced it*; by building new stacks or languages, challengers often just "trade one lock-in for another" rather than solving the core problem of portability. This lack of a truly neutral, performant alternative has consolidated CUDA's position as the only safe bet for large-scale infrastructure.

## Conclusion: Is CUDA the De Facto Standard?
**Yes.**

Based on our findings, *CUDA remains the undisputed de facto standard for HPC and AI*. While SYCL is emerging as a powerful, modern C++ alternative and HIP provides a necessary bridge for AMD hardware, their raw code prevalence is not near that of CUDA. The industry is talking about portability in academic papers, but it is still writing CUDA in the repositories.

To dethrone CUDA, challengers don't just need a better language - they need to replicate twenty years of optimized libraries and developer trust. This is why the industry needs a solution such as [SCALE](https://scale-lang.com) that doesn't ask users to throw away their production-tested investments, but rather **unlocks bringing that source code to new hardware**.

## Appendix

### Case Studies on Performance Portability
* **Bioinformatics (SW# Case Study):** Researchers evaluating the Smith-Waterman protein database search found that SYCL achieves performance comparable to CUDA on NVIDIA hardware while providing "remarkable code portability" to AMD and Intel architectures. 
  * *Source:* [Comparing Performance and Portability between CUDA and SYCL for Protein Database Search (2023)](http://www2.sbc.org.br/sbac/2023/pdfs/sbacpad/s5p2.pdf)
* **The Gaia Mission (AVU-GSR):** A performance assessment of the ESA Gaia mission pipeline demonstrated that HIP reached a 0.94 performance portability score, while SYCL reached 0.93, proving these alternatives can achieve over 90% of native CUDA efficiency in complex astrophysical solvers.
  * *Source:* [Performance Portability Assessment in Gaia (IEEE, 2024)](https://iris.unito.it/retrieve/c1a5f32d-d407-4b14-a076-e878e277c6b8/PP_GAIA_TPDS.pdf)
* **Molecular Dynamics (GROMACS):** Developers adapted GROMACS to use SYCL to target AMD-based supercomputers like Frontier and LUMI, confirming that portability is achievable without major compromises to performance at exascale.
  * *Source:* [GROMACS on AMD GPU-Based HPC Platforms (CUG, 2024)](https://cug.org/proceedings/cug2024_proceedings/includes/files/pap149s2-file1.pdf)
* **Virtual Screening (LiGen on LUMI-G):** A study on the LUMI-G supercomputer showed that SYCL implementations could effectively handle molecular docking on AMD GPUs, even exceeding native throughput in some optimized out-of-kernel configurations.
  * *Source:* [Unlocking Performance Portability on LUMI-G Supercomputer (IWOCL, 2024)](https://re.public.polimi.it/bitstream/11311/1272929/2/IWOCL24___LUMI_G_benchmark%20%281%29.pdf)

## Footnotes
[^1]: (Reuters) [TAIPEI-NVIDIA CEO HUANG: WE NOW HAVE 5 MILLION CUDA DEVELOPERS AROUND THE WORLD](https://www.marketscreener.com/quote/stock/NVIDIA-CORPORATION-57355629/news/TAIPEI-NVIDIA-CEO-HUANG-WE-NOW-HAVE-5-MILLION-CUDA-DEVELOPERS-A--46882472/)
[^2]: [https://blogs.nvidia.com/blog/2-million-registered-developers-breakthroughs](https://blogs.nvidia.com/blog/2-million-registered-developers-breakthroughs)
[^3]: Note: Repository and User detection for OpenCL did not return any result.
[^4]: Note: Repository and User detection for SYCL did not return any result.
[^5]: Approximative average considering the following sources: [carboncredits.com](https://carboncredits.com/nvidia-controls-92-of-the-gpu-market-in-2025-and-reveals-next-gen-ai-supercomputer/), January 6, 2026 (92%), [ainvest.com](https://www.ainvest.com/aime/share/nvidia-holds-92-ai-gpu-market-share-sustainable-dominance-regulatory-risks-bd32e8/), 14 February, 2026 (92%), [finance.yahoo.com](https://finance.yahoo.com/news/nvidia-secures-92-gpu-market-150444612.html), June 6, 2025 (92% including consumer market share), [finance.yahoo.com](https://finance.yahoo.com/news/nvidias-85-gpu-market-share-210500376.html), January 25, 2026 (85% including consumer market share).

---

## Optimizing CUDA Shuffles with SCALE

URL: https://scale-lang.com/posts/2026-01-19-optimizing-cuda-shuffles
Published: 2026-01-19
Author: Justine Khoo

Compilation of CUDA shuffles to use AMD's Data-Parallel Primitives (DPP) is a novel and unique compiler optimization in SCALE, first added in version [1.4.2](https://docs.scale-lang.com/stable/manual/CHANGELOG/#release-142-2025-09-19). It's important to be able to execute shuffles efficiently as they are used in virtually all CUDA kernels. Therefore, this compiler optimization is an example of where SCALE can not only solve the problem of CUDA compatibility with AMD GPUs, but also where SCALE can significantly improve performance of existing CUDA code.

DPP is a hardware feature that's unique to AMD GPUs. AMD GPUs have the ability to execute shuffles more efficiently than any other GPU architecture. Our measurements show that DPP on AMD  is about 10x faster than shuffles in CUDA. Therefore, existing CUDA programs can be sped up by compiling them with SCALE and running them on an AMD GPU, in turn reducing compute costs or increasing the size of model that can be used for a given amount of compute.

However, DPP has been rarely used in practice. This is because no compiler has previously been able to generate code that uses DPP, so DPP could only be accessed via hand-written assembly code or compiler intrinsics. Every program needing to use DPP would need to be hand-optimized specifically for AMD. But not any more: now SCALE enables all existing CUDA programs to execute shuffles efficiently using DPP (including those programs that use shuffles via inline PTX).

## Recap: why shuffles are important in CUDA

CUDA provides functions such as `__shfl` to exchange data within a group of threads according to a specified pattern. For example:

* `__shfl(data, 5, 8)`: read from lane 5

![Broadcast lane 5 shuffle](/assets/shfl_broadcast_lane_5.png)

* `__shfl_up(data, 2, 8)`: read from the lane that's two indexes below

![Shift up by 2 shuffle](/assets/shfl_up_by_2.png)

* `__shfl_xor(data, 4, 8)`: XOR the lane index with 4 to get the source lane

![XOR 4 shuffle](/assets/shfl_xor_4.png)

## The common but slow way of compiling shuffles: ds_bpermute

Upstream LLVM adopts a correct but overly general (and therefore slow) method when compiling the most commonly-used shuffles for AMD GPUs. We will illustrate this with the following example:

**Task**: divide the threads into groups of 16, and share the value in lane 5 with its entire group

**CUDA**: `__shfl_sync(~0u, value, 5, 16)`

![Broadcast lane 5 to group of 16](/assets/broadcast_lane_5_group_16.png)

Here is the generated device code (we'll explain the most important part afterwards):

```
s_load_b128 s[0:3], s[0:1], 0x0
v_mbcnt_lo_u32_b32 v0, -1, 0
v_mov_b32_e32 v2, 0
s_delay_alu instid0(VALU_DEP_2)
v_lshlrev_b32_e32 v0, 2, v0
s_waitcnt lgkmcnt(0)
global_load_b32 v1, v0, s[0:1]
; wave barrier
s_waitcnt vmcnt(0)
ds_bpermute_b32 v1, v2, v1 offset:20
s_waitcnt lgkmcnt(0)
global_store_b32 v0, v1, s[2:3]
s_endpgm
```

The most important instruction for our purposes is `ds_bpermute`. In this context, the word "permute" means to copy data between threads in the warp. `ds_bpermute` can perform an arbitrary permutation across a warp. You simply give it an array `src` of the same size as the wavefront, where `src[i]` is the index that lane `i` will read from. 

`ds_bpermute` is usually more complex than we need, since it can perform any arbitrary permutation between 32 (or 64) lanes. While programmers can make `__shfl` and friends perform any permutation (by passing in a variable as the source lane), the most common cases are those described above: broadcasts, xors and shifts. While the generated code is correct, the generality of `bpermute` comes at a cost: it uses the `__shared__` memory hardware to do the permutation, giving it a comparable cost to a shared memory read.

## Optimizing common types of shuffle: AMD's awesome DPP

The performance of the example above would be significantly improved if the whole computation could be executed while staying within the registers. AMD GPUs have hardware features (including DPP) that can do restricted kinds of permutation an order of magnitude faster. We just need to implement some compiler magic to enable permutations  to be compiled to DPP.

**DPP** is a powerful feature of AMD GPUs which can implement specific patterns of shuffles entirely within the register file. Even better, the permutation can be fused onto the input operand of another instruction, allowing direct access to registers from other lanes when performing many calculations. Using DPP, we've gone from a shuffle needing to access the memory cache, to in many cases being effectively free!

To illustrate the difference between shuffles that need `bpermute` and shuffles that can be executed using DPP, the following diagram shows a complex example of where `bpermute` would be necessary, contrasted with a case that can be optimized with DPP:

![Comparison of bpermute vs DPP shuffles](/assets/bpermute_vs_dpp_comparison.png)

Other patterns supported by DPP include:

* quad perm: do an arbitrary permutation within each group of 4 consecutive lanes.
* row shift: Shift each chunk of 16 right/left by 1-15 lanes. The first/last few lanes are unaffected.
* row rotate: Similar to the above, but with wraparound
* row share: broadcast a specific lane to its entire row
* row mirror: reverse each row
* row half mirror: reverse each octet
* wave64 only:
    * broadcast15: broadcast the last lane of each row to the entire next row
    * broadcast31: broadcast the last lane of the first half to the entire second half
    * wavefront shift: shift the entire wavefront right/left by 1 lane. the first/last lane will read its old value.
    * wavefront rotate: like the above, but the last lane's data goes to the first lane (or vice versa)

We've taught the SCALE compiler to recognise if a shuffle matches any of the DPP patterns and generate the appropriate variant when possible. In cases where no DPP pattern matches,  SCALE will fall back to using `bpermute`.

![DPP operation patterns](/assets/dpp_operation_patterns.png)

## shuffles for free: fusing shuffles into other operations

Since DPP isn't an instruction, it can be fused into other operations. If you want to shuffle some data, then pass the result to another operator, you can do both with just one instruction. 

The reasons that DPP is so much faster than `bpermute` are:

* DPP can be merged with another operation (like an add), whereas `bpermute` cannot. 
* `bpermute` uses the shared memory cache, which will be very busy in real applications that use shuffles. This means that we can expect worse latency from bpermute in real programs that are doing shuffles.

As an example, consider `v_add_u32_dpp v0, v1, v2 quad_perm:[2,3,0,1]`. Here, `v0/v1/v2` are "vector general purpose registers" which are basically registers with a separate value for each thread. Thread 0's `v0` is separate from thread 1's `v0`. This instruction first shuffles the contents of `v2` according to the pattern `quad_perm:[2,3,0,1]`, then adds it with `v1` and stores the result to `v0`. 

![DPP fused with add operation](/assets/dpp_fused_add_operation.png) 

## Optimizing complex shuffles

When a single shuffle primitive is not enough, SCALE combines multiple instructions to build the required permutation efficiently.

### Optimization through sequences of DPP operations

Some common shuffles can't be represented by a single DPP operation, but could be implemented with a sequence of multiple DPP operations that implement the equivalent permutation. Doing a few DPP operations is still more efficient than a single `bpermute`.

For example, the XOR pattern isn't supported on all architectures, but we can build it using a combination of quad perm, row mirror, and row half mirror. Here's how we would do XOR 5:

![XOR 5 implementation sequence](/assets/xor_5_implementation.png)

### Optimizing more complex patterns through swizzle and permlane

AMD GPUs also offer other fast mechanisms for doing permutations that were not commonly exploited prior to SCALE, such as `swizzle` and `permlane`. In general, these mechanisms are a middle ground between `bpermute` and DPP:  faster than `bpermute`, but slower than DPP. They can handle some permutations (for example, a pattern commonly used in FFTs) that are not otherwise representable with DPP.

## Reductions

Two of the most common tasks done with CUDA shuffles are reductions and scans (aka cumulative sums). 

A reduction is a fold over a certain operator over a list of items, like taking the sum of an array. 

![Reduction example](/assets/reduction_example.png)

A scan computes the reduction of all previous lanes for each lane.

![Scan example](/assets/scan_example.png)

CUDA has reduction intrinsics, such as `__reduce_add_sync`. But users can also write their own, using a sequence of shuffles. For example:

```
int acc = in[laneIdx];
for (int offset = 1; offset < warpSize; offset *= 2) {
    acc += __shfl_xor(acc, offset);
}
```

An important insight is that if the user wants to compute a reduction, they can write this in multiple ways. They can do it using what's called a "butterfly reduction", a sequence of xors (like the code snippet above). Or they could do a scan instead, using a sequence of shifts, and extract the value from the last lane. They might even use PTX asm!

![Butterfly reduction with XOR pattern](/assets/butterfly_reduction_xor.png)

![Scan reduction with shift pattern](/assets/scan_reduction_shift.png)

The shuffle optimisations we have been talking about certainly help with reductions and scans, since they are made of shuffles. But we may miss some opportunities to optimise if we only look at each step in isolation. This is because some sequences have individual shuffles that cannot be done by any of the more specialised hardware instructions, only `bpermute`. In the second sequence above, the shifts cannot be lowered to DPP. But if we know the overall result is a reduction, we can generate a completely different sequence of shuffles – where each one can be lowered to an efficient hardware instruction – whose *overall* effect is the same.

So we made the SCALE compiler even smarter: it can now recognise the shape of reductions or scans. Regardless of how the user writes the reduction, the compiler will know what they're trying to do and generate the optimal sequence of machine instructions.

## Conclusion

By optimizing one of the most commonly used instructions in real CUDA code, SCALE can significantly improve performance of existing CUDA code. We are (to our knowledge)  the only compiler that uses DPP operations on AMD GPUs. Upstream LLVM is aware of these instructions, but does not ever generate them!

This topic on shuffles is just one data point of a larger pattern we've observed: AMD GPUs are not constrained by the hardware, but the software. Hence our philosophy at Spectral: you can do amazing things with compilers. AMD's hardware is great and has the ability to achieve high throughput, but the bottleneck is the *software*.

The original developers of CUDA realised the importance of software nearly two decades ago, and the importance of developing software that is able to fully leverage the underlying silicon. This has led to the mature, user-friendly platform that we see in CUDA today. In contrast, until the release of SCALE, AMD GPUs have had so much underutilised potential.

While CPU compiler engineering is a mature field and optimisations have been deeply explored, GPU compiler work is still relatively new. SCALE is full of novel optimisations which (to our knowledge) are not found anywhere else.

---

## GROMACS on SCALE: A Quick Update

URL: https://scale-lang.com/posts/2025-12-16-GROMACS_on_SCALE_December_Update
Published: 2025-12-16
Author: Manos Pavlidakis

After our first GROMACS post [back in May](https://scale-lang.com/posts/2025-05-19-GROMACS), we kept digging into the parts of the CUDA backend that GROMACS relies on, especially after getting a very helpful email from the GROMACS Team. Their feedback pointed out a few places where our original evaluation or implementation wasn’t quite right. So this post is a short update on what changed since then and where things stand now.

## CUDA Graphs: Finally Working

One of the biggest gaps in our previous setup was *CUDA Graph support*. GROMACS uses CUDA Graphs a lot more heavily than we initially realized. Several tests were failing for exactly that reason. Over the past few weeks we implemented the missing pieces:

* Recording events (including the oddities around [`cudaEventRecordWithFlags`](https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__EVENT.html), [`cudaStreamBegin`/`EndCapture`](https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__STREAM.html#group__CUDART__STREAM_1g793d7d4e474388ddfda531603dc34aa3) with the global semantics)
* Making stream capture and graph execution behave the way CUDA does
* Proper handling of graph nodes for kernel launches and copies
* The IPC event edge cases that show up only in real workloads

After wiring all that up, the full GROMACS CUDA test suite now passes on SCALE. This was a good milestone for us because it shows that the CUDA backend inside GROMACS is interacting with SCALE in the way it expects.

## Feature Coverage: SCALE vs HIP, SYCL, and CUDA

Another thing that came up was the comparison with HIP. The HIP backend in the official GROMACS 2025 release doesn’t offload everything to the GPU yet. PME tasks run on the CPU, which affects both performance and comparability. SCALE, on the other hand, supports the same tasks as CUDA tasks because they’re implemented in CUDA inside GROMACS already. Consequently, SCALE can offer a drop in replacement for the CUDA version of GROMACS removing the need to maintain multiple code bases. To verify that SCALE supports the same features as CUDA, HIP, and SYCL we execute the unit tests provided from GROMACS and notice that we execute the same number of tests as all these codebases. 

## Updated Performance Results

We also redid our evaluation using the latest GROMACS release (2025.4) and a more consistent setup based on GROMACS team suggestions:

* HIP is run with `-pme cpu -bonded cpu -update cpu`  
* We use wall-clock time normalized to SCALE as the comparison metric  
* `-resetstep` helped to remove noise in short runs  
* We used [`smi`](https://rocm.docs.amd.com/projects/amdsmi/en/latest/) to inspect GPU utilization instead of relying on indirect counters

A few takeaways:

* SYCL is currently the fastest backend on AMD GPUs. That shouldn’t surprise anyone — SYCL support in GROMACS has been tuned for years and is used heavily on systems like [LUMI](https://lumi-supercomputer.eu/).  
* SCALE is one average 1.5× slower than SYCL but about *1.6× faster than HIP*.  
* GPU utilization of SCALE and SYCL is much higher than HIP - which is because parts of the workload of the HIP case are executing on the CPU.

This aligns well with our goals: for now, SCALE’s priority is compatibility, not beating hand-tuned backends like SYCL. Performance work will come next. The figure below shows the speed-up or slow-down (when values are below one) for SCALE-HIP, and SCALE-SYCL. For our evaluation we use an AMD Radeon PRO MI210 (gfx90a), the GROMACS version is 2025.4, and the ROCm version is 6.3.1. We also tried ROCm 7.1.0 but [AdaptiveCpp does not support it](https://manual.gromacs.org/documentation/current/install-guide/index.html#sycl-gpu-acceleration-for-amd-gpus) and we ended up with compile errors in GROMACS.  
 ![SCALE speedup over SYCL and HIP](/assets/SCALE_speedup_over_SYCL_and_SCALE_speedup_over_HIP.png)

## What’s Next

This round of fixes — especially CUDA Graphs — puts us in a much better position to optimize performance without worrying that we’re missing functional pieces. Our next steps will focus more on kernel code optimizations and other parts of the runtime where we know there’s room to improve. We'll also be testing GROMACS compiled with SCALE to run on NVIDIA GPUs (*sic*).

More updates soon.

---

## It's Been a Minute

URL: https://scale-lang.com/posts/2025-11-10-Announcement_Venture_Backed
Published: 2025-11-10
Author: Michael Søndergaard

## A Big Update from Spectral Compute, the team behind SCALE.

Hey World

Yeah, we know, it’s been a while. We’ve been heads-down, buried in code, and frankly, just incredibly busy making CUDA work *everywhere*.

So, grab a coffee. We’ve got a lot to catch you up on.

We’re [Spectral][spectral], the company behind SCALE, which we’ve been quietly developing for the past six years. Then, last year, we nervously pushed the "release" button on the SCALE beta.
It blew up. Seeing our work hit the front page of [Hacker News][hn] and get covered by [Phoronix][phoronix], [Tomshardware][tomshw] and so many other news outlets was a wild ride.
That response confirmed what we always believed: the world needs SCALE. A way to use CUDA code on any chip. It's why this year, we're all in. We’re now a venture-backed startup, having raised $6 million from our incredible new partners at Costanoa Ventures, Crucible Capital, and 12 amazing angel investors - all to bring SCALE to the entire chip developer ecosystem.

[spectral]: http://spectralcompute.com
[hn]: https://news.ycombinator.com/item?id=40970560
[phoronix]: https://www.phoronix.com/news/SCALE-CUDA-Apps-For-AMD-GPUs
[tomshw]: https://www.tomshardware.com/tech-industry/new-scale-tool-enables-cuda-applications-to-run-on-amd-gpus

## We’re All-In on SCALE

We genuinely believe SCALE is as revolutionary for the GPU space as C compilers were for CPUs.

While it may be harder than creating a new language or trying to JIT what comes out of NVIDIA's compiler, we believe true ahead-of-time compilation of CUDA code is the right solution to this problem, just like it was 20 years ago for bringing C to more CPU architectures. We’re building the entire toolkit to achieve that: we’re implementing CUDA-X library replacements, CUDA Runtime and Driver APIs, etc. And we actually found our toolkit outperformed alternatives, in many cases (average speedup for this test suite: 4.67x)

![SCALE vs HIP - Rodinia Benchmarks - September 2025](/assets/SCALE_vs_HIP_Rodinia_benchmarks_Sept25.png)

## GPU Portability or, a Maze of Dead Ends

The CUDA software ecosystem is massive. Because of that, nobody really wants to reinvent the wheel. To break free from the hardware vendor lock-in, the approach was to port, use translation layers or new languages, all of which have distinct drawbacks.

| Tool | ZLUDA | HIP | OpenCL | Mojo |
|------|-------|-----|--------|------|
| **Benefits** | "CUDA Compatible" | Open source | Open standard | Supports NVIDIA and AMD |
| | Can support (most) programs without recompilation | "Looks like CUDA" | Supported by most vendors | Cool new language |
| **Drawbacks** | Requires use of NVIDIA toolkits | Requires significant effort in porting and subsequent maintenance | Restricted language features | Significant porting and integration efforts from CUDA required |
| | Being rewritten because of legal clawback from AMD | Mainly targets AMD (usually fails on NVIDIA due to target-unlisted backend) | Significantly smaller ecosystem | In making "Python" performant, it ends up looking like CUDA C++ without getting the ecosystem benefits |
| | No AMD data center grade GPU support | Significantly smaller ecosystem | | |
| | Needs DLL injection, flagged as malware on Windows | Significant stability and support issues [^1]| | |
| | Unfinished code has problem* (loss of info needed for re-optimization for AMD architectures) | | | |
| | Cannot match performance of native CUDA | | | |
| | Uses HIP runtime | | | |

That’s why we’re placing all our chips (*eheheh*) on SCALE. It’s our bet - and we think it’s the most clear-sighted approach for the entire ecosystem: abstracting the complexities at the compiler layer and finally run these workloads on any architecture.

## Our Secret Sauce: The Team

We’ve had more than a few people tell us, "*what you're doing is impossible.*". And, you know what? It *is* incredibly difficult. That’s probably why nobody else has been crazy enough to take on this mission. But our R&D team has a combined experience that’s closing in on *two centuries*. If anyone can deliver on this promise, it’s this crew.

We are an amazing, growing team of around 20 people, including some brilliant new folks in our sales and marketing departments (welcome!). But the real MVPs, the absolute strength of this company, are our engineers.

Thanks to them, we’re steadily shipping new versions of SCALE that are already running in production and being tested in PoCs with new prospects as we speak.

## Announcing: SCALE v1.4.2 now with CUTLASS Support

Speaking of shipping... we’re stoked to announce the release of **SCALE v1.4.2**, which (drumroll, please) **includes support for CUTLASS**!

This is a massive milestone for us. We had previously focused on the broader HPC area of compatibility. Now, getting CUTLASS to work at all has been the first giant compatibility hurdle in the AI/ML workload part of the equation, and we've officially cleared it.

Now, to manage expectations (and be totally transparent): it's not winning any speed records just yet. This first release is all about proving compatibility. Right now, our compiler R&D team are deep in the trenches, doing the hard work to optimize and deliver the usable performance you'd expect. You can grab SCALE v1.4.2 [from here][dllink] and use it for non-commercial purposes - tell us what you break.

[dllink]: https://docs.scale-lang.com/stable/manual/tutorials/how-to-install/

This is the necessary groundwork to bring SCALE to the wide world of AI/ML workloads, and it gets us one giant leap closer to...

## What's Next on the Roadmap

**PyTorch support**! This is the big beast we’re currently focused on taming. We have a *long* list of prospects who’ve told us, "*Call us back as soon as you have PyTorch running.*" We hear you, and we’re on it.

Then we need to cover vLLM, SGLang and CuTile, presented by NVIDIA at Pytorch Conference 2025 last month, which we attended. We are genuinely hyped about this. Tools like Triton are cool, no doubt. But CuTile is hitting on something we've been waiting for: bringing that powerful, tiled approach inside the CUDA ecosystem, where it belongs. NVIDIA is set to either open-source the CuTile->TileIR compiler and push it to upstream LLVM, or at least provide open specs for it. Either way, they're (of course) keeping the final machine-code compiler proprietary. That, right there, is our cue. We’ll be the ones building the backend that connects that new paradigm straight to AMD hardware.

Considering all of the above, along with substantial performance enhancements and a few cool secret projects, 2026 is poised to be the year SCALE unlocks the majority of the workloads you’d expect it to happily gulp down and compile. Get ready.

## Come Say Hi!

We’re building all this cool stuff, and now want to make sure people actually know about it We’ve revamped our websites (go check out [spectralcompute.com](spectralcompute.com) and [scale-lang.com](scale-lang.com)) and given our brand a fresh coat of paint.

We’re also getting out in the press more. We just had an awesome interview with Business Insider. You can check it out [here](https://www.businessinsider.com/spectral-compute-funding-pitch-deck-nvidia-cuda-2025-11).

We’ll also be at SC25, the Conference for High Performance Computing in St. Louis.
We are so excited to be there to talk about... well, all of this. Come say hi at Booth #6552.

If you want to get serious and talk shop, you can schedule a 1-on-1 time slot with us at this [link](https://calendar.app.google/KUxJhdpqZxkRc65g8).
Otherwise, just swing by the booth! We’d love to have a chat, and you’re more than welcome to "steal" as much of our cool new handouts as your bag can hold. We’re also planning some chill, post-event meet-ups. If you want to grab a beer with the team, drop us a line on our [Discord][discord] or reply to the [Reddit post on the topic][reddit-post].

[discord]: https://scale-lang.com/s/discord?utm_source=blog
[reddit-post]: https://www.reddit.com/r/CUDAUnlocked/comments/1othfdr/some_big_news_from_the_team_we_got_funding_and/

That’s the update. It feels good to be back. We’ll try not to be strangers.

Cheers, yours

- Michael.

[^1]: [MI300X vs H100 vs H200 Benchmark Part 1: Training – CUDA Moat Still Alive – SemiAnalysis](https://semianalysis.com/2024/12/22/mi300x-vs-h100-vs-h200-benchmark-part-1-training/)

---

## Running GPU-Optimised Monte Carlo (GOMC) on an AMD GPU using SCALE

URL: https://scale-lang.com/posts/2025-06-16-GOMC
Published: 2025-06-16
Author: ftvkyo

SCALE is a toolkit that enables CUDA applications to be natively compiled for AMD GPUs.
As SCALE reaches maturity, the number of projects that work with it steadily increases.
Today we will look at the experience of using SCALE to build and run an existing CUDA-enabled project called [GOMC][gomc].

[gomc]: https://gomc-wsu.org/

You can read more about SCALE in the [official documentation][scale-docs] or in the form of a blog post: [Unpacking SCALE][scale-post-unpacking].
However, reading those pages is not necessary for understanding this post.

[scale-docs]: https://docs.scale-lang.com/stable/
[scale-post-unpacking]: /posts/2024-08-02-unpacking-SCALE

## What is GOMC?

GOMC is short for GPU-Optimized Monte Carlo.
This open-source software allows simulating molecular systems using statistical methods.
It can run on CPUs using OpenMP and on NVIDIA GPUs using NVIDIA CUDA.
You can find more information about GOMC on its official [website][gomc].

This software is optimised for running on the GPUs, but as it uses NVIDIA CUDA, by design, it cannot be used with GPUs from other vendors.
As such, GOMC is a perfect candidate for using SCALE that removes that artificial barrier.

## Installing SCALE

Before we can move on to GOMC, we need to install SCALE.
I will use SCALE Free Edition version 1.3.1, which is the most recent version at the moment of writing.

My GPU is *AMD Radeon RX 7700 XT 12GB GDDR6*, its Instruction Set Architecture (ISA) is `gfx1101`.
If you want to replicate my results, please make sure that your GPU is on the list of [GPUs supported by SCALE 1.3.1][scale-supported-gpus].

[scale-supported-gpus]: https://docs.scale-lang.com/1.3.1/#which-gpus-are-supported

You can find the table of ISA names in [AMD docs on GPU specs][amd-archs].
Look for the column "LLVM target name", even though these identifiers are not specific to LLVM.
SCALE also includes a way to [determine your GPU ISA][scale-determine-isa] once you install it:

```sh
scale-free-1.3.1-Linux/bin/scaleinfo | grep gfx
```

[scale-determine-isa]: https://docs.scale-lang.com/1.3.1/manual/how-to-use/#identifying-gpu-target
[amd-archs]: https://rocm.docs.amd.com/en/docs-6.3.1/reference/gpu-arch-specs.html

**Remember:** academic/research licenses are avaliable, and even if your GPU is not on the list, or is only supported by the Enterprise Edition,
[get in touch][scale-contact-us], and we will try to help you.

[scale-contact-us]: https://docs.scale-lang.com/1.3.1/#contact-us

SCALE can be installed by following the official instructions: [Install SCALE 1.3.1][scale-install].
I am using neither Ubuntu, nor Rocky Linux, so I will follow installation instructions for "Other Distros":

[scale-install]: https://docs.scale-lang.com/1.3.1/manual/how-to-install/

```sh
# Download the tarball
wget https://pkgs.scale-lang.com/tar/scale-free-1.3.1-amd64.tar.xz

# Extract it into the current directory
tar xf scale-free-1.3.1-amd64.tar.xz
```

A `scale-free-1.3.1-Linux` directory will appear.
This marks the end of SCALE installation process.

## Buliding GOMC with SCALE

The GOMC build process described here is based on SCALE's validation scripts.
You can learn more about them in their repository: [spectral-compute/scale-validation][scale-validation].
We also mentioned them in the [post about GROMACS][scale-post-gromacs].
You can find the GOMC validation script I used as a reference here: [scale-validation/gomc @ dc1f32c][scale-validation-gomc].

[scale-validation]: https://github.com/spectral-compute/scale-validation
[scale-post-gromacs]: /posts/2025-05-19-GROMACS
[scale-validation-gomc]: https://github.com/spectral-compute/scale-validation/tree/dc1f32c84a6ec9932544094bc5f1a40d66fce777/gomc

Before we start building GOMC, we should activate [SCALE virtual environment][scale-env] (`scaleenv`).
The concept is similar to virtual environments used with programming languages like Python.
While using `scaleenv` is not strictly necessary for using SCALE, it greatly streamlines the experience:

- `scaleenv` simplifies configuring and building projects by managing common environment variables automatically
- `scaleenv` makes dealing with multiple SCALE versions and/or multiple GPUs easier

[scale-env]: https://docs.scale-lang.com/1.3.1/manual/how-to-use/#the-easy-way-scaleenv

To activate the virtual environment, you need to specify which GPU ISA you want to use:

```sh
# Activate the virtual environment for gfx1101 in the current shell.
# Substitute "gfx1101" with your GPU ISA.
source scale-free-1.3.1-Linux/bin/scaleenv gfx1101
```

Your shell prompt will be modified to reflect that the virtual environment is active.
Should you wish to deactivate the virtual environment, simply run `deactivate`.

Now, let's clone the sources of GOMC:

```sh
git clone https://github.com/GOMC-WSU/GOMC.git && cd GOMC

# Check out a specific revision of GOMC.
# This is different from the revision in the validation script,
# as I simply used the newest one available.
git checkout 9fc85fb574194e93752776c44c00c3e6513af281
```

As can be seen from the validation scripts, GOMC currently requires a few small patches to make it work.
The following three patches fix what can be considered bugs in GOMC:

```sh
# Don't enable libc++.
# Reason: CMake detects that we are Clang, GOMC enables libc++ for that in a way that seems to be broken.
sed -i -E 's/-stdlib=libc\+\+//' \
    "CMakeLists.txt" \
    "lib/cereal-1.3.0/CMakeLists.txt"

# Disable OpenMP support.
# Reason: GOMC's OpenMP support tests __GNUC__'s value.
# Even if a larger value is used that causes use of the newer OpenMP syntax that Clang expects, it still results in error messages.
sed -i -E 's/find_package\(OpenMP\)//' "CMakeLists.txt"

# Disambiguate a call to Difference.
# Reason: this is also required when using NVIDIA CUDA.
# See this ticket: https://github.com/GOMC-WSU/GOMC/issues/492
sed -i -E 's/ Difference/ ::Difference/g' "src/GPU/CalculateForceCUDAKernel.cu"
```

The following patch works around a genuine compatibility issue SCALE will need to fix in a future release.
GOMC hardcodes `CUDA_ARCHITECTURES "35;60;70;80"`, and this patch removes `35` from that list.

```sh
sed -i -E "s/(\s+CUDA_ARCHITECTURES ).*/\1\"60;70;80\"/" \
    "CMake/GOMCCUDASetup.cmake" \
    "test/BuildGPUTests.cmake"
```

That patch was necessary because GOMC provides a definition for `double atomicAdd(double *, double)` when `__CUDA_ARCH__ < 600` to compensate for the lack of that API in older NVIDIA CUDA versions.
On the other hand, in SCALE 1.3.1, `atomicAdd` for `double` exists for values of `__CUDA_ARCH__` below that.

This illustrates how much effort goes into making a compatible API.
We can't just provide the necessary functions, we have to provide *no more than necessary*.
Otherwise, function and type declarations start clashing, [SFINAE][sfinae] breaks and so on.
We publish [API Coverage Reports][scale-api-coverage] in our documentation.
We also list [differences and enhancements][scale-differences] in SCALE compared to NVIDIA CUDA.

[sfinae]: https://en.wikipedia.org/wiki/Substitution_failure_is_not_an_error
[scale-api-coverage]: https://docs.scale-lang.com/1.3.1/manual/apis/
[scale-differences]: https://docs.scale-lang.com/1.3.1/manual/differences/

GOMC can now be configured and built.
GOMC instructions advise using `./metamake.sh`, but we found it easier to interface with CMake directly:

```sh
cmake \
    -DCMAKE_BUILD_TYPE=RelWithDebInfo \
    -DGOMC_GTEST=On \
    -B"build" \
    .

cd build
make -j$(nproc)
```

At this point, GOMC has been built.
You can verify that it works by running some tests:

```sh
./GOMC_GPU_GCMC_Test --gtest_filter="-ConsistentTrajectoryTest.*"
./GOMC_GPU_GEMC_Test --gtest_filter="-ConsistentTrajectoryTest.*"
./GOMC_GPU_NPT_Test --gtest_filter="-ConsistentTrajectoryTest.*"
./GOMC_GPU_NVT_Test --gtest_filter="-ConsistentTrajectoryTest.*"
```

## Running GOMC with SCALE

Let's run an example provided by GOMC:

```sh
# Assuming we are still in the GOMC build directory, save the path to it for later
export GOMC=$(pwd)

# Leave the GOMC repository
cd ../..

# Clone GOMC Examples and check out a specific revision.
# Like with GOMC itself, I simply chose the most recent one.
git clone https://github.com/GOMC-WSU/GOMC_Examples.git && cd GOMC_Examples
git checkout c3d5d2e2c52d98db2d143019e8fe36f8fe0e62cc

# Select the example we are going to run
cd NVT_GEMC/pure_fluid/octane_T_360_00_K

# Run the example
"$GOMC/GOMC_GPU_GEMC" in.conf 2>&1 | tee test.log
```

The output will contain information about the GPU being used:

```
GPU information:
Info: Device Number: 0
Info: Device name: AMD Radeon RX 7700 XT - gfx1101 (AMD) <amdgcn-amd-amdhsa--gfx1101>
Info: Memory Clock Rate (KHz): 1124000
Info: Memory Bus Width (bits): 192
Info: Peak Memory Bandwidth (GB/s): 53.952000
```

This is what simulation results look like for me:

```
Simulation Time (total): 4.7654e+02 sec.

Move Type                Mol. Kind      BOX_0      BOX_1
% Accepted Displacement        C8P     51.75844   74.14317
% Accepted Rotation            C8P     50.76357   86.17870
% Accepted Regrowth            C8P     21.82463   59.14750
% Accepted Crank-Shaft         C8P     84.44648   86.98939
% Accepted Mol-Transfer        C8P      0.26120    0.46801
% Accepted Volume-Transfer             54.22542   54.22542

Displacement:                           51.6984    sec.
Rotation:                               10.6198    sec.
MultiParticle:                           0.0000    sec.
Brownian-like MultiParticle:             0.0000    sec.
Intra Molecule Transfer:                 0.0000    sec.
Regrowth:                               44.1079    sec.
Intra Molecule Exchange:                 0.0000    sec.
Crankshaft:                             29.8462    sec.
Intra Targeted Transfer:                 0.0000    sec.
Molecule Exchange:                       0.0000    sec.
Molecule Transfer:                     176.1490    sec.
Nonequilibrium Molecule Transfer:        0.0000    sec.
Targeted Transfer:                       0.0000    sec.
Volume Transfer:                       135.4983    sec.
```

I was able to take some screenshots of the inputs and outputs of the simulation using [VMD software][vmd].

[vmd]: http://www.ks.uiuc.edu/Research/vmd/

Input file `STEP3_START_C8P_sys_BOX_0.pdb`:

![](/assets/posts/gomc-input-box-0.png)

Input file `STEP3_START_C8P_sys_BOX_1.pdb`:

![](/assets/posts/gomc-input-box-1.png)

Output file `octane_360_00_K_BOX_0.pdb`:

![](/assets/posts/gomc-output-box-0.png)

Output file `octane_360_00_K_BOX_1.pdb`:

![](/assets/posts/gomc-output-box-1.png)

## Wrapping up

As one of the developers of SCALE, I am happy to see it work this well.
There are many projects that are built exclusively for NVIDIA CUDA or that support GPUs from other vendors only as an afterthought.
By bridging that gap, SCALE, among other things, makes research work more accessible, and I find that important.

I chose GOMC as the subject of this post specifically because I wanted to produce some visual results.
In SCALE development, we deal with a lot of text, such as unit test outputs, stack traces, build system logs.
Because of that, I thought it would be exciting to produce some pictures and share them here :)

Overall, I would say that the experience of using GOMC with SCALE has been pleasant.
The patches that I had to apply were minor, and 3 of them were justified by the fact that some assumptions about the environment that GOMC made were no longer correct.
The experience of using `scaleenv` was nice, and from my experience, that was even easier than dealing with NVIDIA CUDA or AMD HIP.

There are, as always, things that could be improved.
If you'd like to submit a bug report, here is how you can do it: [Report a Bug][scale-bug].
In the meantime, [join our Discord][scale-discord] and check out our [F.A.Q.][scale-faq]

[scale-bug]: https://docs.scale-lang.com/stable/contact/report-a-bug/
[scale-discord]: https://scale-lang.com/s/discord?utm_source=blog
[scale-faq]: https://docs.scale-lang.com/stable/manual/faq/

## Acknowledgements

In the making of this post, I used the following:

### GOMC (GPU-Optimised Monte Carlo)

- MIT Licence
- [Website](https://gomc-wsu.org/)
- [GitHub Repository](https://github.com/GOMC-WSU/GOMC)

### GOMC Examples

- MIT Licence
- [GitHub Repository](https://github.com/GOMC-WSU/GOMC_Examples)

### VMD (Visual Molecular Dynamics)

- Custom Licence
- [Website](http://www.ks.uiuc.edu/Research/vmd/)

> VMD was developed by the Theoretical and Computational Biophysics Group in
the Beckman Institute for Advanced Science and Technology at the University of
Illinois at Urbana-Champaign.

> Humphrey, W., Dalke, A. and Schulten, K., `VMD -Visual Molecular
Dynamics', J. Molecular Graphics, 1996, vol. 14, pp. 33-38.

VMD is recommended by the [GOMC manual](https://gomc-wsu.org/Manual/software_tools.html#vmd).

### SCALE Free Edition version 1.3.1

- [Custom Licence](https://docs.scale-lang.com/1.3.1/licensing/)
- [Website](https://scale-lang.com)
- [Third-party software used by SCALE](https://docs.scale-lang.com/1.3.1/notices/)

---

## SCALE Benchmark case study: GROMACS

URL: https://scale-lang.com/posts/2025-05-19-GROMACS
Published: 2025-05-19
Author: Chris Kitching

With the release of version 1.3.1, SCALE has reached a major compatibility milestone: the ability to run the CUDA version of GROMACS on AMD GPUs.

This is an exciting moment for us, since GROMACS is a large and complex CUDA project. Getting to this point involved overcoming a number of stubborn issues in our compiler, and is a positive sign for how mature the SCALE platform is becoming.

Since version 2025.0, GROMACS has begun to offer AMD GPU support via a HIP port. Our work is unrelated to that effort: SCALE compiles CUDA source code directly to AMD machine code, no porting required. This allows us to run even older versions of GROMACS that predate the HIP port.

We've tested correctness on GROMACS 2024.4 and 2025.1, and are excited to present the results of our performance benchmarking on 2025.1.

## Performance Comparison to HIP port

Since GROMACS has an ongoing HIP port, it offers a valuable opportunity to perform a direct performance comparison. Projects that have no existing AMD GPU support are more challenging to evaluate, since it will necessarily involve comparison of runs on NVIDIA hardware against those on AMD hardware. While valuable, such runs are not ideal as a direct measurement of the performance impact of SCALE: we want to measure how close we’re getting to our goal of matching or exceeding the performance of a HIP port.

The chart below shows the relative performance of HIP (ROCM 6.3.1) vs. SCALE 1.3.1 on the Max Planck Institute GROMACS benchmarks, running on a single MI210 GPU. Negative performance improvement represents cases where SCALE was slower than HIP, while positive values represent performance improvements.

![](/assets/gromacs-comparison.png)

As shown in the graph, our performance is broadly comparable to HIP, and sometimes significantly faster. An important caveat here is that the HIP port is incomplete, so some of these wins are simply a result of SCALE using the GPU for more of the work. To reduce this effect, we disabled the CUDA acceleration of PME tasks for the SCALE runs, since those currently lack a HIP port. As a result, real-world users of GROMACS with SCALE should see somewhat better performance than depicted here.

To get a more fine-grained understanding of the differences seen here, we turn to GROMACS’ internal performance counters. These measure the time spent in each part of the application, allowing us to directly compare the performance of individual kernels (or groups thereof) compiled with SCALE vs. their HIP replacements. Since there is not a 1-1 mapping between CUDA kernels and HIP kernels in the two implementations, this approach is more meaningful than simply timing individual kernels.

The output from these performance counters is shown in the table below:

![](/assets/gromacs-table.png)

The cases where SCALE is dramatically ahead (e.g. “Force”) are examples of cases where the HIP code simply hasn’t been written yet, so it’s comparing CUDA-with-SCALE to CPU, which isn’t particularly interesting. What’s encouraging is that in the other cases, SCALE is still coming out slightly ahead or on-par with HIP, suggesting that when GPU code does exist in both versions, the results are broadly similar on each. These improvements represent cases where our compiler or runtime library has outperformed the HIP version.

These measurements demonstrate the idea that CUDA code can be compiled for AMD hardware without a “performance tax": it’s just a matter of doing the necessary compiler/platform work to make it happen. As we iterate on performance in the coming months, we anticipate more exciting results to share with the community. Since SCALE is a platform, not a per-project porting effort, performance benefits will benefit a wide range of CUDA applications, not just the ones for which we publish benchmarks.

## Try it yourself

If you want to reproduce these results yourself, you can find the scripts we use to build, test, and benchmark GROMACS here: [https://github.com/spectral-compute/scale-validation/tree/master/gromacs](https://github.com/spectral-compute/scale-validation/tree/master/gromacs)

That repository contains similar scripts for other projects we use for validation. Check out the README for more information, and consider contributing scripts for projects you care about. Adding real user code to SCALE’s test suite helps us improve the product further.

The free edition of SCALE offers support for RDNA GPUs. If you want to obtain evaluation/academic access to SCALE Enterprise \- which offers CDNA support \- please [get in touch](mailto:hello@spectralcompute.co.uk)

Bugs? [Send a report](https://docs.scale-lang.com/stable/contact/report-a-bug/)

## Looking forward

Although the results presented above are encouraging, we recognise that there's still significant work to do. So far, our focus has been on expanding SCALE's API completeness/compatibility to the point where enough real-world projects work out of the box. For our next major release, we are switching focus more towards improving performance of the platform. We believe that SCALE is significantly more compelling when it has competitive (or superior) performance compared to HIP: if the CUDA version of your project is just as fast, then using SCALE offers a promising alternative to a costly HIP port.

---

## August in review

URL: https://scale-lang.com/posts/2024-08-28-August-in-review
Published: 2024-08-28
Author: Michael Søndergaard

It’s been about six weeks since we released SCALE and we’re overjoyed with the reception we’ve received. We:

* Made number 1 on Hacker News and stayed there for several days!
* Saw more than 35 articles written about SCALE.
* Had more than half a million views across various videos made about SCALE.
* Welcomed over 250 people to the SCALE Discord

A successful launch, wouldn’t you say? 

It seems several companies agree with us, as we’ve been contacted by over a dozen organizations for commercial purposes & partnerships, and all that while we’re still in beta! We weren’t quite expecting this level of interest but are happy to report that talks are underway and we’re looking forward to sharing more once the details have been ironed out.

Speaking of beta, we’re hard at work getting the next patch for SCALE out and you can expect it to drop very soon.  

In the same vein, the consulting side of our business has been booming as well. We do a little something we like to internally call ‘’cost-imization’’. In reality this is performance/compute optimization, but the end goal is typically to reduce costs and it has a nice ring to it. 

Sometimes, this means that we optimize the software running an AI without changing the behavior of the model itself to massively reduce cloud compute costs. 
Other times, we’re reworking older software to resolve longstanding technical debt, freeing up large amounts of compute to ‘’do the thing’’. 
Although cost-imization isn’t always the goal, sometimes it’s just about speed. In those cases where every millisecond counts, we’re squeezing every last bit of performance that we can out of the software, bringing latencies down as low as they can go.

We enjoy doing this kind of work because our team is primarily composed of performance specialists and GPU experts with a lot of experience in developing low latency software - although that’s not really doing the team justice but going into that would be a blog post all on its own. Suffice to say that the intertwined and synergistic specialties of our developers are the whole reason we set out and built SCALE in the first place: we had the right people for the job.

And that’s it for today! As mentioned above, we have a patch incoming next week and you can expect an accompanying post along with that. Once again, we’d like to thank everyone for their continued excitement, bug reports and support in making SCALE a success!

---

## Unpacking SCALE

URL: https://scale-lang.com/posts/2024-08-02-unpacking-SCALE
Published: 2024-08-02
Author: Matthew Ireland

It’s been just over two weeks since the first beta version of SCALE launched. We’re thrilled by the response we’ve received: it’s clear that the world is as excited as we are at the prospect of diversifying the GPGPU market\!

SCALE enables CUDA code, previously restricted to NVIDIA GPUs, to be run on AMD GPUs. This is a massive step forward in the GPGPU programming world: it opens up a huge number of GPGPU projects to AMD customers who previously could not access them. Even more importantly, SCALE enables end-users to make a free choice of GPU hardware (for example, to opt for powerful but lower-cost GPUs from AMD). As more vendors enter the GPGPU market over the next decade, SCALE will provide a unified cross-vendor platform for software developers which will motivate GPU vendors to compete based on their hardware offering, rather than relying on software lock-in to achieve market dominance.

## What is SCALE?

SCALE is a GPGPU programming toolkit that enables CUDA applications to be natively compiled for AMD GPUs. Before SCALE, programs written in the CUDA programming language could only be compiled for NVIDIA hardware, using tools provided by NVIDIA. Using SCALE, CUDA programs can be run, unmodified,  on AMD GPUs. The great thing about SCALE’s approach is that it doesn’t add any overhead, allowing the performance of CUDA code on AMD GPUs to equal or exceed the performance of the same code on an NVIDIA GPU.

## What SCALE **is not**

SCALE is:

1. Not a new language that has incompatibilities with CUDA.

1. Not a “translation layer”. It doesn’t use NVIDIA runtime libraries in any way. 

1. Not based on NVIDIA tools in any way (it was built in a cleanroom environment).

1. Not an emulator. There’s no inherent slowdown when using SCALE.

1. Not based on HIP, or any other framework that’s incompatible with CUDA. We’ve written our own CUDA-compatible framework from the ground up.

SCALE is not affiliated with AMD, NVIDIA, or any other hardware vendor. SCALE has been built by Spectral Compute, a small UK company, whose staff are truly committed to diversifying and evolving the GPGPU market.

## Why did we build SCALE?

CUDA code dominates the GPGPU landscape, and it’s clear why: CUDA is a great language\! Millions of lines of CUDA code have been written in thousands of projects across applications such as deep learning, AI, bioinformatics, computational fluid dynamics, weather forecasting, and many more. Large numbers of organisations have invested years of engineering effort into building large CUDA codebases, and built up significant expertise in the CUDA language.

Now, if any project or organisation with a significant CUDA codebase wishes to support not just NVIDIA GPUs, but also AMD and other vendors, they would have had a significant problem before the release of SCALE. They would have had several options:

1. Manually rewrite their entire code-base in a different language that can be compiled for both NVIDIA and AMD GPUs, in particular using open standards such as SYCL. However, this is a huge undertaking and could take many man years for a large codebase. Subtle bugs would likely remain after the translation effort is complete, and would take further time to resolve.

   Adopting an open standard would be a viable approach for new projects (in fact, we’re huge fans of SYCL\!). However, the overhead of translating an existing large codebase to a new language is simply not viable for many organisations.

   There’s also the issue that vendors will never be incentivised to support open standards to the extent that they support their own implementations. Use of API wrappers would remain necessary. Vendors are likely to always focus their efforts on their own ecosystems.

1. Translate their codebase into a new language such as HIP, which comes with an automated tool for translating from CUDA. This can work well, but HIP is not intended to be a drop-in replacement for CUDA. Therefore, this approach can still incur significant overheads in terms of extensive manual tuning and bug fixing.

   The CUDA code that has been rewritten using HIP can also no longer be compiled using NVIDIA’s own tools, and so must rely on wrapper APIs. For this reason, many projects end up maintaining separate CUDA and HIP codebases.

1. Use a JIT compiler such as ZLUDA. There’s a key advantage to this approach: since ZLUDA operates on PTX assembly code, an end-user can run a CUDA project on an AMD GPU without having access to the original source code, and without support from the project developers. While JIT necessarily incurs a modest slowdown compared to a natively-compiled program, the performance is likely to be sufficient in practice for most applications.

   However, code run using ZLUDA can never have the same performance as binaries that have been natively compiled and opportunities for target-specific optimisation are limited. Furthermore, ZLUDA relies on DLL injection, which can be fragile, and tends to interact badly with antivirus software. Careful attention must also be given to NVIDIA’s EULAs, which are unfavourable towards the use of translation layers.

   ZLUDA is not currently under active development. This is unfortunate, since it does fill a useful niche.

1. Write your own CUDA emulator. However, this would be a huge undertaking and incur a slowdown. This would defeat the entire purpose of GPGPU programming: performance.

We set out to give any organisation with a CUDA codebase a real choice of hardware platform, without needing to rewrite any code, and without the migration introducing any new bugs into their software. In this way, SCALE was born: a toolkit that allows CUDA applications to be natively compiled for, and automatically optimised for, AMD GPUs.

SCALE is a unified cross-vendor platform that will constitute the future of GPGPU programming. CUDA programmers will be able to devote their time to adding new features to their codebases, rather than rewriting them for cross-vendor support. The market will be diversified: new vendors will be able to enter the market and existing CUDA code will *just work* on their platforms. Existing vendors will be forced to compete based on the value their hardware actually brings to end-users, rather than achieving dominance through software lock-in.

## Is SCALE really 100% CUDA-compatible?

Yes. We even built [support for inline PTX](https://docs.scale-lang.com/manual/inline-ptx/). But there’s some nuance: there are multiple dialects of CUDA, and we built SCALE without any reference to the NVIDIA implementations (as required by the NVIDIA licence agreements). It’s a requirement to:

1. Achieve bug-for-bug compatibility with the dialect of CUDA used, and

1. Support all the undocumented features of the dialect of CUDA used, as well as the documented ones.

This might sound complicated. However, it’s actually quite straightforward: we offer two compilers (one for each of the two major CUDA dialects) and our automated test framework includes all the tests for a large number of third-party projects (in addition to a huge number of our own tests\!). This enables us to understand what semantics people expect from CUDA code in the real world, and make sure that SCALE provides those semantics.

In a way, whenever someone writes a unit test in an open-source CUDA project, they’re also writing a unit test for SCALE. We benefit from a huge number of third-party unit tests that allow us to ensure the correctness of SCALE, without relying on any NVIDIA tools.

However, SCALE is still in beta status: there are currently some known defects in SCALE where we do not yet achieve perfect compatibility with CUDA. We have taken care to [document these defects](https://docs.scale-lang.com/manual/differences/), and are working hard to get them fixed. There are also a number of missing CUDA APIs and features, which we’re continuing to add to SCALE.

## And there’s really no slowdown when using SCALE?

There’s no inherent reason that code compiled with SCALE should be slower than code compiled with NVIDIA’s toolchain. On the contrary: SCALE has the potential to outperform NVIDIA’s toolset.

However, it is important to remember that AMD hardware is not NVIDIA hardware. There’ll be some things that AMD can do better, and some things that NVIDIA can do better. This is the great thing about SCALE facilitating cross-vendor support: it will drive forward performance of the underlying hardware as vendors are forced to compete based on the features they can offer in silicon.

Lots of CUDA code has been heavily-optimised for NVIDIA GPUs (and, in fact, for specific NVIDIA targets). Therefore, it’s likely that some performance tuning will be necessary when migrating to AMD GPUs, just as tuning would be necessary when migrating to a newer model of NVIDIA GPU.

Code compiled using SCALE will *work correctly* on an AMD GPU straight away. Performance tuning is a continuous process, as new generations of hardware evolve, and will never stop. Adding a few \#ifdefs to your code will always be cheaper than a complete rewrite.

## Further advantages of SCALE

There’s a further advantage to our independence from NVIDIA and other GPU vendors: we can add new features to CUDA, while still maintaining 100% compatibility with it. We’re passionate about enabling programming productivity through the use of features that are expected in a modern programming language, and using static analysis to find as many problems in code as possible at compile time.

We’ve [already announced](https://docs.scale-lang.com/manual/runtime-extensions/) a few of the features that SCALE adds to CUDA (including the option to use exceptions instead of C-style return codes), and there are many more to come in the near future.

## Conclusions

Even in its current beta status, SCALE facilitates the diversification of the GPGPU market by enabling CUDA code to be compiled for AMD GPUs. This enables AMD customers to use a huge number of projects that were previously inaccessible to them, and enables everyone to benefit from a choice of GPU platform based on the underlying performance of the hardware.

SCALE is not an emulator or a translation layer. Therefore the performance of SCALE can match and even has the potential to exceed the performance of code compiled using NVIDIA’s own toolchain. SCALE is a true cross-platform solution and maintains independence from NVIDIA, AMD and other hardware vendors.

We see a future where organisations can speed up their codebases by picking the vendor that provides the best hardware features for them. We see a future where organisations can support their users regardless of which hardware platform their users opt for. We see a future where vendors focus on novel hardware developments, and new vendors can break into the GPGPU market. And we’re really excited about this future\!

---

## Announcing the SCALE BETA

URL: https://scale-lang.com/posts/2024-07-12-release-announcement
Published: 2024-07-12
Author: Michael Søndergaard

Announcing SCALE: a GPGPU toolchain allowing CUDA programs to be natively run on AMD GPUs. Support for more vendors is in development.

We want everyone to have the freedom to use the best GPGPU programming tools, and the GPU hardware that best suits their needs. Write once, run anywhere.

[Get started](https://docs.scale-lang.com) with SCALE today.

## Why SCALE?

We believe that it should be possible to write code once, and build/run it on any hardware platform. This has been a reality for CPU code for many years, so why not GPUs?

We set out to directly solve this problem by bridging the compatibility gap between the popular CUDA programming language and other hardware vendors.

One codebase, multiple vendors.

## What is SCALE?

SCALE is a GPGPU toolkit, similar to NVIDIA's CUDA Toolkit, with the capability to produce binaries for non-NVIDIA GPUs when compiling CUDA code.

SCALE aims to be source-compatible with CUDA, including support for inline PTX and nvcc's C++ quirks.

Instead of providing a new programming language for GPUs, SCALE embraces the widely-popular CUDA language and allows its use with non-NVIDIA hardware. This frees developers from the challenges associated with maintaining multiple versions of their codebase.

## Beyond CUDA

SCALE supports a *superset* of the CUDA language. It also offers various opt-in [language extensions](https://docs.scale-lang.com/manual/language-extensions/) that go beyond the featureset of NVIDIA's CUDA, improving developer experience and productivity. Why should multi-vendor support be _less_ convenient than a single-vendor solution?

[Download the free SCALE beta](https://scale-lang.com) today, and begin the liberation of the GPUGPU market.

## Dive Deeper

Jump into [the technical manual](https://docs.scale-lang.com) for all the gory details.