The what

xc is a C-like language I've (where "I" includes Claude) been building over the last 6-7 months. Its compiler, xcc, automatically turns data-parallel code-candidates into executable blocks across various workers {CPU threads, CUDA, Metal, Vulkan or WebGPU} from the same source. On NVIDIA it generates PTX itself and loads it through the driver, so a program needs no toolkit, nvcc or extra libraries. All it needs is the driver that ships with the card.

The same source code compiles for all targets. It runs on Metal on an M4 Max and on CUDA on Windows, unchanged.

The specific block (CPU, GPU, ...) that is executed by the scheduler is bound at runtime, not compile-time, based on a measurement of the performance on the actual hardware present. Once the measurement has been taken, it is persisted to the standard 'settings' location on the platform ('defaults' on Mac, Registry on Windows, ~/.config on linux,...). If a code-block is to be run with sufficiently-different parameters than before, the measurement is re-made, with the goal being to establish the boundary condition between CPU and GPU. In this way it self-optimises.

Crucially, it's the same language used to create both CPU and GPU (and in future HLS to FPGA) code-blocks. The compiler then has the ability to prove (at the IR stage) that there are no data-hazards, and to interpose transport/memory-barriers as required to enforce that.

Recently, I've added a "C front end", the CFE uses clang to produce LLVM IR, and converts their IR into my IR and then everything flows as normal. That means plain, unadorned C code that has loops that can be proven to be non-self-dependent can be parallelised onto the GPU.

I say "GPU" but in fact xcc provides what I call 'parallel parallelism' - the built-in scheduler will quite happily tile work across multiple GPUs on a given machine, or in fact over a network (preferably high speed) or if you're on a Mac, the Thunderbolt 4/5 connections. The repo provides 'xcd' - a daemon that is run on the remote machines, which announces itself using mDNS and can take part via the scheduler.

The scheduler measures, for each block

• Its transfer time (to and from the worker), and
• Its compute time, at tile sizes that double from a minimum (×2, ×4, ×8, …).

It uses the smallest tile that reaches 90% of the best throughput, so the work isn't drowned by the transfers. For a networked worker, the smallest tile is set from the link's round trip and bandwidth, keeping latency under 5% of a tile's time.

'auto' (the default unless the programmer wants to force a particular worker for a particular block of code) then times the devices alone and together, and keeps the fastest arrangement. Within a shared run, workers take tiles from a common queue, so a faster worker naturally does more of the work. On my windows box with a 3090 and a 5080, the 5080 gets roughly 2x the work of the 3090.

The example

There are several benchmark-style results on the website, CPU, GPU, the CFE, and networked. but I wanted to see it in a real-life scenario rather than just artificial benchmarks. The application I chose was one that is pretty painful to run on a CPU - I call it Astronomancy, and it's an astrophotography frame stacker.

xc provides a cross-platform UI framework (though this is in its infancy), which is useful when you have a cross-platform compiler :) Astronomancy leverages that framework to be runnable on any of the three main platforms (Windows, Mac, Linux) without source changes. The below image is on a Mac, but the same works under either of the other two.

To generate those beautiful astrophotography images, the camera (usually a high-definition one) takes many many photos over a period of time, and the software needs to merge them together, take account of inherent camera noise, debayering, stuck pixels, passing satellites, etc. etc. The pipeline is

• calibration
• star detection and registration
• Lanczos-3 warping and Bayer drizzle
• sigma-clipped and median stacking
• background extraction
• wavelet denoising
• Richardson-Lucy deconvolution with deringing
• star separation
• stretching with local contrast

Every per-pixel stage is a par block, and none of it is hand-written CUDA

I took a short sequence (640 or so, each about 12MB) of photos of M51 (the whirlpool galaxy) together with the appropriate darks etc frames, and ran it through a processing chain...

Here's an example of what the program is working with, a single frame from the source data, and the final result after merging

M51 data: Erik Westermann (formerly RemoteAstrophotography.com), Vespera/Stellina, used with credit.

The timings, when run in various modes, are interesting:

You can tell from the numbers that there's a lot of code which is 'saxpy'-like, where the overhead in moving the data to the GPU is a significant fraction of the compute time involved. Astronomancy is written in xc, which allows annotation of the 126 'par' blocks where data-parallel execution would be possible - if it were written in C, the CFE would find them instead.

In 'auto' mode, for each of those blocks, it chose the most efficent way to compute the block of code. Putting all the loops etc. on a single GPU was not as fast as sharing between both GPUs on the machine, but the real win came from deciding which of the candidate blocks would be best-served by putting on the GPU, bringing the time down from a naive single-threaded version taking almost half an hour to approximately 2 minutes.

[aside: The Mac has an advantage on one of the blocks being (frequently) used that needs 48GB/s of bandwidth for an image, maxing out the PC's 5950X's dual-channel DDR4-3200 but the M4 Max has more than 10x that bandwidth and copes much better, the actual compute-time on the PC GPU is faster than the Mac, but it's dominated by the memory-bandwidth requirements, specifically it's the weford op:

• Items: 3072 × 2080 × 3 planes = 19.2 M.
• Bytes per item: 6 reads + 3 writes × 4 bytes = 36.
• Per run: 36 × 19.2 M ≈ 691 MB.
• On the PC: 691 MB in 14.4 ms = 48 GB/s, against a 51 GB/s peak, so it's at the limit.]

I have a policy of using real-life programs to test things, and (as usual) it turned out to be useful..

• Stale GPU copies. A block writing through stack + k*n marked only the slice's start as changed. CUDA then reused a stale copy of the whole array on the card, so the master dark came out all zeros. Vulkan and Metal kept no such copy, so they were right. Fixed in the 0.87.1 release
• A race that crashed the app on Windows. CPU threads wrote into an array that a GPU block had just written and was still tracking, which crashed with 0xC0000005. Fixed in the 0.87.1 release.

Documentation is in the first link at the top, and also contains the repo link.

The compiler is only at v0.87, theres quite a way to go before I want to call it 'version 1' but its already fairly capable. There are certainly places it could be improved (arrays of double currently keep a block on the CPU under CUDA, for example, because it doesn't work across all platforms. I need to relax some things to allow host-specific operations to work).

Let me know if you have questions :)

tileiras contains only 5 new passes and 30 MLIR dialects

It appears that CuTe DSL compiler based on the same LLVM fork as cicc (54 new passes) and so produces fastest code, also it supports 32 MLIR dialects and scripting from python

Hi,

I am an embedded software engineer who's work experience is on bare metal microcontrollers and embedded Linux systems. I have mainly worked on MCUs such as STM32, IoT devices such as ESP32, etc. using the skills which a lot of embedded folks would be familiar with - C/C++ programming, I2C, UART, SPI comms, sensor interfacing, signal processing, interrupt handling, debugging registers etc.

I did a Masters degree recently in Computer Engineering and took courses on systems programming, that involved more :

OS skills such as multithreading, synchronization and concurrency, memory management and data structures, and a

Linux kernel programming course that involved Linux processes, threads, virtual memory, and scheduling, Memory allocation, page faults, and NUMA, CPU affinity, context switching, and synchronization, Shell scripting, GDB, perf, and system tracing etc.

Computer architecture course that got into cache organization and locality, pipelined architectures, performance optimizations, etc.

I'm seeing a lot of roles around GPU and ML infrastructure, and wanted to know how realistic would it be to transition to such roles from where I am. Is this knowledge helpful or do I have to know more about machine learning for that?

Hello, I have good understanding of the GPU primitives. I am trying to write the custom kernels for my engine. As of now I am directly using flashinfer etc. so I want to know how did you optimize your kernel for your use case?

2x T4 on Kaggle, no P2P between the cards, 70W cap each. Running Qwen3.8-27B with tensor parallel (exllamav3, NCCL backend) plus a DFlash2 draft model. Decode is ~55 tok/s overall, but every verify step pays an all-reduce over PCIe through host memory.

Questions for people who've done this on PCIe-only cards: 1. Has anyone overlapped the all-reduce with the next layer's GEMM on T4 (sm75) and got a real win? Streams, chunked all-reduce, anything. 2. Is a custom all-reduce through pinned host memory ever faster than NCCL when P2P is off? 3. The cards sit at 450-1000 MHz under the power cap. Does it change what's worth optimizing?

Can share nsys traces if useful.

So I am in in the middle of bringing my app out, and trying to stabilize pipelines, however for image generation, VULKAN runs 4-6x faster than cuda and I can't figure out why?

Any help is much appreciated.

We're running a hackathon to build a simplified PTX compiler running from October 21st to October 27th, 2026.

The goal is to generate efficient PTX from an internal kernel representation. The parser is provided, so you don't need to implement the full CUDA compiler frontend, and you don't need to find an efficient schedule. The goal is to generate efficient PTX code to run the program.

Here is an example of IR for RMSNorm:

    Tile[a0, a1_co] (N=8192)
        for a1 in a1_co..1024:128
            in2 = load x[a0, a1]
            f32 v1 = multiply(in2, in2)
            f32 acc0 <- add(acc0, v1)
        WarpShuffle(acc0, length=32)
        Smem float acc0_smem[4]
        if ((lane == 0)):
            acc0_smem[warp] = acc0
        Sync
        TreeHalve(acc0, length=4, tid=warp)
        in1 = load mean_count[0]
        f32 v0 = reciprocal(in1)
        f32 v2 = multiply(acc0, v0)
        in0 = load add_c1[0]
        f32 v3 = add(in0, v2)
        f32 v4 = sqrt(v3)
        f32 v5 = reciprocal(v4)
        for a2 in a1_co..1024:128
            in3 = load x[a0, a2]
            f32 v6 = multiply(in3, v5)
            in4 = load w[a2]
            f32 v7 = multiply(in4, v6)
            mul_1[a0, a2] = v7    Tile[a0, a1_co] (N=8192)
        for a1 in a1_co..1024:128
            in2 = load x[a0, a1]
            f32 v1 = multiply(in2, in2)
            f32 acc0 <- add(acc0, v1)
        WarpShuffle(acc0, length=32)
        Smem float acc0_smem[4]
        if ((lane == 0)):
            acc0_smem[warp] = acc0
        Sync
        TreeHalve(acc0, length=4, tid=warp)
        in1 = load mean_count[0]
        f32 v0 = reciprocal(in1)
        f32 v2 = multiply(acc0, v0)
        in0 = load add_c1[0]
        f32 v3 = add(in0, v2)
        f32 v4 = sqrt(v3)
        f32 v5 = reciprocal(v4)
        for a2 in a1_co..1024:128
            in3 = load x[a0, a2]
            f32 v6 = multiply(in3, v5)
            in4 = load w[a2]
            f32 v7 = multiply(in4, v6)
            mul_1[a0, a2] = v7

The prize pool is $ 2,750, and we offer internship opportunities to participants with interesting submissions.

Info: https://riftstack.ai/hackathon

Intro lectures:

• YouTube: https://youtu.be/KK6TWf4Pp6A?si=qXlYK2jop4K68lqZ
• Substack: https://kernelspace.substack.com/p/building-a-gpu-compiler-from-scratch

21 new instructions:

• applypriority.async.bulk (1 forms)
• applypriority.async.bulk.tensor (12 forms)
• clmad.hi (2 forms)
• clmad.lo (2 forms)
• fabric.submit (1 forms)
• fabric.try_atom.async (13 forms)
• fabric.try_get.async (2 forms)
• fabric.try_get.async.tensor (10 forms)
• fabric.try_pullred.async (5 forms)
• fabric.try_put.async (8 forms)
• fabric.try_put.async.tensor (20 forms)
• fabric.try_red.async (16 forms)
• fabric.try_red.async.tensor (20 forms)
• fabric.wait (1 forms)
• mbarrier.check_layout (1 forms)
• multimem.red.async (1 forms)
• multimem.st.async (3 forms)
• spcompress (1 forms)
• spdecompress (1 forms)
• tcgen05.ld.red.spcompress (1 forms)
• tcgen05.ld.spcompress (1 forms)

I'm debugging GPU P2P on a Minisforum MS-01 (i9-12900H):

• Internal RTX 2000 Ada 16 GB on PCIe.
• External RTX A2000 12 GB through an AOOSTAR USB4 dock.
• NVIDIA 595.99.02, BIOS 1.26.

Stock CUDA does not expose peer access. With experimental driver/libcuda modifications, A2000→Ada PUSH passes, while Ada→A2000 PUSH and both PULL directions fail.

This persists under native Linux without Xen, with corrected synchronization and full payload comparisons. CUPTI confirms separate SM and copy-engine tests. Physical CE copies also retain the asymmetry.

Could this reflect an Ampere/Ada restriction, requester behavior, or a bug in our experimental mapping path? Which NVIDIA driver code path or diagnostic would you examine next?

I cannot swap the cards or put them behind one switch. Sources, logs and negative results:
https://github.com/PieroBorgatta/CaskLogic-PCIe-P2P-Lab

Thanks for the help

Hi Everyone,

I wanted to follow up regarding my previous post and ask how realistic would it be to work in ML systems and inference roles, if my professor's main expertise is in vector quantization and recommendation systems, similarity search etc. Can I ask him if he's open to working on accelerating GPU inference for ML or upskill on these things myself. Could you suggest some topics where I can leverage the area he works in that can use inference and performance optimization? Can I get an opportunity to work on jobs such as these -

1
2
3
4

Thank you

Hi all, I previously share GPUFlight here, a CUDA profiler I'm building. I've now added Rust support to its Performance Lab and Workbench.

After seeing NVIDIA’s CUDA Rust announcement, I wanted to try writing CUDA kernels in Rust using cuda-oxide. I was also curious whether GPUFlight's gpufl trace command could profile them. It runs your CUDA application and collects profiling data that you can inspect in the GPUFlight UI.

It worked after I fixed a few issues, so I added Rust support to Performance Lab. You can now solve the lab problems in Rust and profile your kernels. GPUFlight’s Workbench also supports Rust if you want to experiment with your own code.

I've marked it as Beta because I'm still working on benchmark calibration, but it's ready to try.

Try Performance Lab

If you'd like to learn more about "gpufl trace", check out this guide to profiling a CUDA application: Profile a CUDA app with gpufl trace

seems that ptxas has bug in its register allocator and marks uniform registers for desc operand as dirty - and so allocating and assign new pair every time

I got offered a PhD position in parallel computing (in the US). But given the way AI is taking over fields such as math and software engineering, do you think fields such as parallel and distributed computing, GPU programming etc. will have the same fate in the next few years? Will the value of having specializations go down? Thanks.

Hey everyone,

For the past few months, we've been building **ACE-Silix**—a specialized kernel language and compiler toolchain (`.slx`) designed to compile fused AI operations directly to silicon with zero intermediate DRAM staging, targeted specifically at heavily constrained edge hardware.

Rather than relying on generic high-level dispatchers, Silix compiles kernels down to native hardware-bound binaries (protected via NIST AES-256-GCM) and binds them directly into a static, single-replay CUDA Graph.

We just completed our full end-to-end audit on an **NVIDIA Quadro T1000 (4GB VRAM, Turing SM75, 0 Tensor Cores)** as shown in the benchmark trace above.

### Key Takeaways:

1. **Memory-bound operations show massive gains:** Removing intermediate DRAM staging in RMSNorm and RoPE yielded a **4.30× speedup (69.8 µs vs 300.3 µs)** and **1.78× speedup** with zero floating-point divergence (`1.00000000` Cosine Similarity).

2. **Amdahl’s Law inside captured graphs:** While isolated MLP and AdaLN kernels run faster, full graph execution converges to compute-bound parity (~44.34 ms vs 44.46 ms) with slightly tighter variance (±0.28 ms).

3. **DDIM Scheduling Regression:** The DDIM step showed a 13.1% regression, which we traced to register pressure and suboptimal unrolling in the current compiler pass—actively refining this.

We are currently packaging a drop-in ComfyUI node (for DiT acceleration) and a lightweight local LLM runtime binary for edge testing.

Curious to hear from systems engineers and CUDA devs: what bottlenecks do you run into most frequently when deploying DiT or LLM architectures to memory-constrained hardware once graph launch overhead is already eliminated?

I measured fp16 vs fp32 energy on a Tesla T4: 8192×8192 matmul ×100, 5 runs each, alternating, under three settings.

Setting fp16 fp32 GPU energy time

No cooling 475 J · 8.6 s 2.53 kJ · 39.7 s -81.2% -78.3%

Start ≤65 °C 426 J · 7.8 s 2.21 kJ · 33.8 s -80.8% -77.0%

≤65 °C + 3 warm-ups 413 J · 8.4 s 2.50 kJ · 40.1 s -83.5% -79.0%

All p = 0.008 (Mann-Whitney). fp16 was 4.3–4.8x faster.

What I found useful:

- Warm-ups plus a steady start temperature brought the run-to-run spread down to 0.1–0.2%.

- fp32 sat at the 70 W power cap 92% of the time (fp16 ~60%), so part of the gap is the card's limit. A fixed-power-limit run is next.

- Energy is from NVML (GPU only), not wall power.

Tool: pip install landauer-gap, then

joules compare --cool-to 65 --warmup 3 "python a.py" "python b.py"

Disclosure: I built it. How do you control for power caps in your benchmarks?

I wanted a clean number for how much GPU energy fp16 saves over fp32 on a T4, so I measured it carefully instead of guessing from TDP.

**Workload:** 8192×8192 matmul, 100 iterations, PyTorch, Kaggle T4.

**Method:** fp16 and fp32 runs alternate (so thermal/clock drift hits both sides equally), 5 runs each, median reported, Mann-Whitney test between the two. GPU energy comes from the NVML energy counter (GPU only, not wall power).

| Setup | fp16 | fp32 | GPU energy saved | Time saved | Spread (IQR) |

|---|---|---|---|---|---|

| No cooling between runs | 475 J, 8.6 s | 2.53 kJ, 39.7 s | 81.2% | 78.3% | 1.0–1.3% |

| Every run starts at ≤65 °C | 426 J, 7.8 s | 2.21 kJ, 33.8 s | 80.8% | 77.0% | 2.6–5.1% |

| ≤65 °C + 3 warm-up runs | 413 J, 8.4 s | 2.50 kJ, 40.1 s | 83.5% | 79.0% | 0.1–0.2% |

p = 0.008 in all three setups (the smallest p possible with 5 vs 5).

**Why the gap is this large:** the T4 is Turing, so it has no TF32. fp16 runs on the tensor cores while fp32 runs on the regular CUDA cores. fp16 finishes 4.3–4.8× faster, and since energy ≈ power × time, most of the saving comes from the shorter runtime.

**Caveat:** the T4 has a 70 W power cap. fp32 sat at the cap 92% of the time and fp16 about 60%, so these runs are power-limited. On an A100/H100 with TF32 enabled I'd expect a much smaller gap.

**Reproduce it:**

```

pip install landauer-gap

joules compare --names fp16,fp32 --baseline fp32 --runs 5 --warmup 3 --cool-to 65 \

"python bench.py fp16" "python bench.py fp32"

```

bench.py is basically:

```python

import sys, torch

dt = torch.float16 if sys.argv[1] == "fp16" else torch.float32

a = torch.randn(8192, 8192, device="cuda", dtype=dt); b = torch.randn_like(a)

for _ in range(100): a @ b

torch.cuda.synchronize()

```

Disclosure: I built joules (the CLI). It's open source and works without an account.

If anyone has A100/H100/L4 numbers for the same test, especially TF32 on vs off, I'd love to see them.

A small detail that matters when first learning CUDA: the CPU can return from a kernel launch before the GPU has finished the work.

A timer around the launch alone may therefore measure submission overhead, not time to a usable result. A complete explanation needs to follow the input buffer, the kernel’s thread indices, ordering and completion, and the result’s next consumer.

For example, four blocks of 256 threads provide 1,024 thread positions. If the input contains 1,000 elements and each thread handles one element, the final 24 positions need a bounds check. Launch geometry and useful data size are related, but they are not identical.

I’m building Systems Unpacked and prepared a beginner explanation of this path. AI assisted the drafting. The examples are conceptual, not GPU benchmark results. Full explanation: https://systemsunpacked.dev/blog/what-is-cuda

For people teaching CUDA: which part most often causes confusion—mapping threads to data, or reasoning about when the work is complete?

Hello! Crossposting this here because it involves CUDA programming as much as it does reinforcement learning. Porting RocketSim to run many thousands of envs in parallel on the GPU was difficult but I learned a lot about GPU hardware & software optimization during this project. It's been in my mind for close to 2 years and has existed in many different half-complete iterations since then but now I have something I'm proud to share. Still rounding out the repo, implementing tests and additional functionality, but wanted to open this up to anyone who's interested :D

I've been working on a runtime for a pattern that is awkward in CUDA: a million independent samples (ODE solves, Monte Carlo paths) where each one stops at a different step. The obvious kernel is one thread per sample with a loop, but threads that finish early keep their warp busy until the slowest thread in it is done.

What eagle does instead:

• it removes finished samples and packs the live ones together between batches of steps;
• it records the loop as CUDA graphs;
• it chooses between a fused multi-step kernel and a persistent kernel based on the batch size, and only times both when the answer isn't clear.

You write the per-sample code once, in typed Python (hawk), and it is compiled at run time (with NVRTC, or nvcc when it is installed).

A million samples in FP64 (Runge-Kutta 4 integration), each stopping at its own step, against a Warp per-thread kernel:

• H100: 613 µs against 1.54 ms (2.5×).
• Tesla T4, the free Colab GPU: 46 ms against 194 ms (4.2×).
• Across the six GPUs I measured: 2.5× to 4.2×. PyTorch with torch.compile is 9.8× to 130× slower.

When all samples run every step, eagle and Warp tie, with Warp slightly ahead. Compared with the GPU's theoretical FP64 limit for this workload, it reaches 46% on the H100 and 65 to 92% on the other five GPUs I measured, so there is room left. I'd love feedback on the persistent kernel and the reordering design. A developer guide and an internals walk-through are the next docs work; until then, questions and issues are welcome.

Personal project; I work at ESA on GPU trajectory software. Benchmark files, scripts and a kit to rerun them: https://amasat01.github.io

(code: github.com/amasat01/eagle and github.com/amasat01/hawk).

I am currently working on step 6.

1. DocumentSASS approach of getting SASS right out of nvcc

2. Reverse engineering of an example cubin (measuring_stick) using nvdisasm

3. Open-ended Angluin's L* over the nvcc + RE derived device SASS forms (hw validation) using a query protocol

4. Build the device SASS rulesets using repeated step 3 queries + hw facts (op chain latency, architecture errata)

5. Build vendor container (empty cubin)

6. Inject transpiled SASS into empty container (your c, cu, cpp gets turned into a relational IR and emitted using the device SASS ruleset derived over 3 and 4.

The goal is 6. context reuse (native SASS streaming)

My latency pass measures part latency directly so stalls and nops are tuned per part, not per arch. Baseline stalls for an architecture are conservatively tuned by nvidia, this query protocol elimitates binning performance gaps, and unlocks the maximum per chip performance by giving you a custom stall profiler that doesn't crash your machine timing out wdt.

Seems that cicc has 42 new passes (not presented in vanilla clang)