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)

Sharing a debugging story because the trace taught me something about reading Nsight that I would been getting wrong.

Setup: Wan 2.1 T2V 1.3B, full fwd + bwd + AdamW step, bs 1, bf16, synthetic inputs, MIG slice (2g.48gb) of an RTX PRO 6000 on a rented GPU pod.

The problem: changing only the AdamW implementation moved my forward + loss time from ~207 ms to ~128 ms (CUDA events, medians over 70 steps, no profiler attached). The optimizer runs after backward, so that shouldn't happen.

The trace for the slow run:

Two things in there I had to learn:

1. The empty regions on the GPU stream (~35 ms and ~69 ms) sit right at the start of a step, inside my forward window. CUDA events measure elapsed time between two markers on the stream, so idle time counts. My "forward time" was partly the GPU waiting.
2. The 71.6 ms cudaLaunchKernel bar is host wall time, not GPU time. The kernel it launched ran for ~7 µs. A long API bar means the host thread was stuck somewhere, not that the GPU was busy.

What the optimizer did: forced single-tensor AdamW is ~6,600 launches per step vs ~560 for foreach and 78 for fused. My best explanation is that with thousands of tiny launches the CPU never gets far ahead of the GPU, so any host stall turns straight into a gap. With foreach and fused the big gaps disappeared, even though foreach's kernels took more GPU time.

Then the host side:

nproc: 256
cpu.max: 1360000 100000

256 visible CPUs, but a cgroup budget of 13.6 CPUs worth of time per 100 ms. The throttling counters in cpu.stat went up during the profiling run, and both time gaps are shorter than the 100 ms period, which fits. I didn't capture scheduler events around these exact steps though, so I can't say for sure throttling caused these particular gaps.

I didn’t set that limit. It came with the rented pod, the same thing a Kubernetes CPU limit would do. Any container with a CPU limit works this way.

What would be interesting to checknext time:

- does the timer include idle time?
- in Nsight, is the GPU slow or just empty?
- what's in cpu.max, not just nproc (especially on rented containers) ?
- capture context switches for the launch thread, so a stall like this can be pinned on the scheduler or ruled out

Longer write-up: 256 visible CPUs, yet my GPU was waiting.

I built a GPU compiler and wrote an article about it, with a focus on the concepts shared by modern compilers like TVM, Modular, and PyTorch Inductor.

The article breaks CUDA kernel generation into a sequence of passes and shows how a common operation (RMSNorm, Matmul, Attention) is progressively lowered to generate a CUDA kernel. Each pass introduces its own IR that is optimized for specific tasks like kernel fusion or schedule enumeration.

Each example is reproducible, so you can write PyTorch code and generate code at a specific IR level, which makes the flow easier to understand.

% emmy compile -c 'nn.Linear(3, 2, bias=False)(torch.randn(4, 3))' --ir loop 
for a0 in 0..4         
  for a1 in 0..2             
    for a2 in 0..3                 
      in0 = load input[a0, a2]
      in1 = load linear_wt[a2, a1]
      v0 = multiply(in0, in1)
      acc0 <- add(acc0, v0)
   linear[a0, a1] = acc0

On the benchmarked workloads, many outperform Torch Inductor and cuBLAS. Generated kernels use modern features like tensor cores and TMA.

The field is very dynamic, with many competing approaches, so it's hard to pin down the "core concepts"; this is my take on them.

YouTube: https://youtu.be/KK6TWf4Pp6A?si=0hCuOMEan4mauJjN

Article: https://aiadvances.org/building-a-gpu-compiler-from-scratch-76d69e98ed03

GitHub: https://github.com/cloudrift-ai/emmy

• cicc does not have MLIR
• cicc has 870 builtins vs 715 in nvptx
• cicc has 3183 intrinsics vs 1998 in nvptx

I've been building this in my spare time after getting tired of spending the first hour of every GPU rental on setup.

How it works: pick a GPU, choose "Data Science", and JupyterLab opens with PyTorch and your saved Kaggle / Hugging Face tokens already in place. In my tests it's been ready in 14–32 seconds. Billing is prepaid and per minute, and unused time is refunded when you stop early.

I'd appreciate honest opinions on:

1. Would you pay a bit more than the cheapest raw GPU for this kind of convenience?
2. What would make you trust a new GPU service with your data and money?
3. Is storage that keeps your files between sessions the obvious next feature?