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.

I ran into a CUDA memory-accounting trap while measuring a hand-built GPT-2 KV cache in PyTorch on an NVIDIA A40, and wrote it up in my first blog post.

The cache formula was 72 KiB per token, but requested_bytes.all.current minus my baseline was consistently 8,519,688 bytes above it, regardless of sequence length. Splitting that offset gave 8,519,680 bytes for the lazily allocated cuBLAS workspace plus 8 bytes for my live int64 next-token tensor.

In a fresh process, I allocated a 64x64 float32 CUDA tensor, recorded requested bytes, ran y = x @ x, then subtracted the 16,384-byte output from the increase. The remaining increase was exactly 8,519,680 bytes on this setup. Before any CUDA allocation, that statistics key was absent, so I used .get(key, 0) for the initial reading.

Running a warmup forward pass BEFORE the baseline, releasing its outputs, and subtracting the remaining non-cache tensors made requested minus formula exactly 0 bytes at every checkpoint from 64 to 1,024 tokens. memory_allocated still had a separate allocator hump around 384-640 tokens, peaking around 9.5 MiB above the requested payload, with smaller bumps later.

My takeaway: warmup matters for memory measurements too, and requested payload is not the same as allocator block sizes. I am not assuming this workspace size holds for every GPU, stream configuration or library version.

For people who profile CUDA allocations: what other persistent library or per-stream allocations do you explicitly warm up or exclude before measuring an operation's incremental memory?

Full experiment and Figure 6: https://angry-kratos.github.io/posts/2026-10-kv-cache/

Code: https://github.com/angry-kratos/angry-kratos.github.io/tree/main/code/kv-cache

Disclosure: this is my blog; AI assistance was used for editing and tooling.

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?

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 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?

Just curious about what are the different applications where you guys are using CUDA? When I mean using CUDA, I mean at the low level, and understanding it's internals, modifying kernels etc and not just calling libraries. What are you guys using it for?

We just open-sourced KAI Pichu, an agent harness for writing CUDA kernels and fusing existing ones. You describe what the operator computes and how accurate it needs to be. No existing kernel or PyTorch implementation required.

You can also give it a chain of kernels to fuse. It generates and optimizes fused candidates, measuring the complete sequence against your unfused baseline.

It compiles, checks correctness, profiles with Nsight Compute, and iterates. You choose the model, GPU, and budget. The code and results from every attempt are available to inspect. There's a bundled skill to help Claude Code or Codex set up your task.

On a 7×7 depthwise conv, we got 3.5–3.9× over cuDNN across held-out shapes on an RTX 5070. The run took about two hours and about $1 in API calls. Full run.

If you've got a kernel chain you want fused, give it a try. I'd like to hear what works and where it gets stuck.

Code and setup

• nvcc code 15–18% faster
• nvcc code is bigger and more amenable to stall counts reducing
• all mlir based compilers use vanilla llvm nvptx backend :-(

I've been experimenting with a different way of explaining Tensor Core programming.

https://svgfpsaiagents.github.io/svgfpsgameAI/mmatensorbankconflict.svg

https://svgfpsaiagents.github.io/svgfpsgameAI/mma_tile_engine.svg

https://svgfpsaiagents.github.io/svgfpsgameAI/hopfphotonics.html

https://en.wikipedia.org/wiki/Light_cone

Instead of writing CUDA kernels or building a heavyweight simulator, I wanted to see how much of the Tensor Core programming model could be represented as a single self-contained executable artefact.

The result is a browser-based model written entirely as one SVG + JavaScript file.

It models:

• GEMM as tiled execution (D = A·B + C)
• tile scheduling (ti, tj, tk)
• accumulator evolution
• FP16 inputs with FP32 accumulation
• warp fragment ownership
• mma-style fragment mapping
• global → shared → register → tensor-core data flow
• shared-memory bank behaviour
• row-major vs padded vs XOR-swizzled layouts
• reuse and memory-traffic estimates
• FP16 vs FP32 numerical error maps
• interactive 3D visualisation of the output matrix

The interesting thing for me is that all of these views are driven from the same underlying state. The bank conflict view, fragment view, tile execution view, error map and result visualisation are all observing the same model as it executes.

The goal wasn't to build a cycle-accurate simulator. It deliberately does not model things like occupancy, scheduling, pipeline hazards or cache behaviour.

Instead, I was exploring a question:

Can Tensor Core behaviour be computationally modelled, rather than merely described, in a compact interpreted artefact that requires no CUDA compiler or GPU while still preserving correspondence between matrix maths, tiled execution, fragment distribution, memory layout, data movement and numerical effects?

One unexpected outcome is that the entire thing remains inspectable. There are no dependencies, no build system and no generated files. The complete model lives in a single source file.

I'm curious whether people here would consider this:

1. a visualiser,
2. an educational simulator,
3. an executable specification of the Tensor Core programming model,
4. or something else entirely.

I'd especially appreciate feedback from anyone familiar with CUTLASS, CuTe, mma.sync, ldmatrix, Tensor Core fragment layouts or GPU architecture education. - Aston Walker EdgeMafia SVG

Hello all,

I've been going through the Programming Massively Parallel Processors (5th edition) book and doing the kernel exercises in CUDA C++. Once I finish that, I wanna start writing CUDA kernels for H100 GPUs (specifically stuff like Flash Attention 3, RoPE, etc). Based on a bit of research, it seems that H100 GPUs have a lot of features not covered in the PMPP book, so I was wondering where I could find one or more solid resources that can get me started with writing programs? So far I've tried searching for some stuff, but they all seem to either not really cover in enough detail or are way too long. Any advice would be much appreciated. Thank you!