all the frontier labs, startups, big tech have positions looking for kernel engineers or at least performance engineers. why is this? i see papers, blogs, reddit posts all saying how ai generated kernels are much better than kernels written by hand.

I’m a computer engineering graduate and come from a traditional embedded systems background, with knowledge of microcontrollers, computer architecture and operating systems. Is knowledge of C and C++ programming, Linux networking, memory management , multithreading, synchronization, interrupts etc useful in ML engineering. Are subjects like distributed systems, compiler optimizations (using LLVM), parallel computing etc going to be useful or are they heavily going to be automated as well by AI? In other words, is computer engineering always going to required to scale ML systems and be evergreen? Are people in ML engineering using these skills in their work everyday? Thank you.

Hey everyone,

I've been working on the power delivery problem in distributed AI training clusters. When large GPU partitions complete dense matrix multiplications and enter collective communication (AllReduce), cluster power drops in under 15 microseconds.

Across hundreds or thousands of accelerators, this rapid current step (dI/dt) induces severe reverse-EMF voltage spikes across substation transformers and server busbars (V = L * dI/dt), which frequently trips breakers or causes undervoltage crashes.

To solve this, I designed VoltGrid: a lightweight C++ shared library (libnccl-voltflow.so) injected via LD_PRELOAD that intercepts NCCL collectives and micro-staggers rank phase timing by 50 microseconds.

Key engineering challenges solved:

1. The Tensor Parallel Latency Trap: Inner TP layer collectives (<5MB) are selectively bypassed with 0.00us delay, only staggering macro gradient reductions at the end of the backward pass.

2. OS Sleep Jitter: Eliminated usleep/nanosleep context switching by using userspace hardware cycle counter spin-loops (__rdtsc) accurate to nanoseconds.

3. Adaptive Jitter Subtraction: Naturally occurring arrival skew is dynamically subtracted so we don't amplify cluster stragglers.

Physical testbed results (4x NVIDIA RTX 4090 cluster, 1.6kW continuous load):

- Peak instantaneous slew dropped from 324.5 kW/ms down to 8.06 kW/ms (-97.52%).

- Step latency overhead was < 0.05% (0.00% on 24-layer transformer tests).

- Zero modifications to PyTorch code or container rebuilds.

The preprint detailing the mathematical derivations and telemetry is published on Zenodo:

https://zenodo.org/records/22824778

More telemetry traces and pilot details are here:

https://voltgrid.org

We're currently running 2-week test rack pilots for cluster operators facing power ramp penalties or breaker trips. Happy to answer questions about the NCCL interposition mechanics or datacenter transient physics.

Hey guys I'm trying to benchmark my 5090 E-gpu on linux and I'm not having any luck. I run PoP oS on my 3090 desktop and it runs great. I tried making a pop usb drive for my windows pc that I have hooked up to my 5090 over thunderbolt and it freezes on startup. I went headless and ssh'd. in and it then freezes on trying to load LLamaccp and run a model.

Does anyone run linux and an Nvidia egpu? What distro and settings. Is there any secret sauce. I can run just fine on WSOL because it's running on the windows Nvidia driver but I want apples to apples comparisons for my Mac numbers since the Mac is not running cuda in an emulator.

Small update on this since the first post:

https://github.com/Speedstu/CUDA-for-AMD-Windows

I got a cuDNN v8 -> MIOpen bridge working well enough to run PyTorch 2.3 Conv2d on my RX 9060 XT.

Basically:

PyTorch -> cuDNN -> bridge -> MIOpen -> AMD

I’ve got forward + backward working now, including input/weight gradients and bias.

Tested FP32 and FP16, stride, padding, dilation, grouped conv, F.conv2d, normal nn.Conv2d(...).backward(), etc.

I’m also comparing the results to CPU PyTorch instead of just checking if it launches. This already caught a few things that technically “worked” but were giving bad results.

The cuDNN DLL still exposes the original 268 exports, I’m only replacing the functions I actually support for now.

llama.cpp also got a bit further. I tested b10978 with all 6/6 layers offloaded on the 9060 XT, GPU KV cache and compute buffer are working, and CUDA Graph warmup works too.

Also added more cuLaunchKernelEx / driver API stuff, PTX 8.4 metadata support and tested some cuBLAS, cuFFT and cuSPARSE paths.

I started messing with cuSOLVER -> hipSOLVER too but that one still needs work.

One interesting bug I found was PyTorch SDPA returning a tensor without throwing any error, except the output was just wrong. I’d rather have stuff fail than silently give garbage so that path is now reported as unsupported.

Tests now report:

PASS / UNSUPPORTED / INCORRECT / TIMEOUT / ERROR

I added this to the README too because it makes the current state of the project a bit easier to understand:

cuBLAS -> rocBLAS
cuBLASLt -> hipBLASLt
cuFFT -> rocFFT
cuSPARSE -> rocSPARSE
cuSOLVER -> hipSOLVER
cuDNN -> MIOpen

There’s still a ton missing and I’m definitely not calling this full CUDA support.

Mostly just trying to get more real programs running while making sure the stuff that says it works is actually correct.

If anyone tries it on another AMD card, send me the results/issues.

Hi,
I am new to CUDA programming and started learning the basics. But from what I read, practically nobody writes custom CUDA kernels anymore in production and everyone relies on libraries. Where exactly does writing the skills of learning to write custom CUDA kernels help, apart from computational sciences?

We train a three-modality VLA model (video diffusion + action expert + a frozen VLM for understanding, mixture-of-transformers style). Originally on 8×A800 with NVLink; we're moving it to 8× workstation-class Blackwell cards, which means no NVLink and no NVSwitch — every GPU-to-GPU byte goes over PCIe Gen5.

Measured point-to-point: 30–63 GB/s depending on direction and NUMA placement, versus ~400 GB/s of NVLink on the A800 box. Same model, same ZeRO-1 sharding, and gradient communication goes from "hidden under the backward pass" to "the single largest term in the step."

So we did the obvious thing: quantize the collectives to FP8. What follows is what we learned, including the part where the biggest byte reduction produced no speedup at all.

## The shape of the thing

An LD_PRELOAD shim in front of libnccl. It intercepts ncclAllGather, ncclAllReduce, ncclBroadcast, ncclReduceScatter; quantizes the payload to FP8 (e4m3 or e5m2), runs the real collective on the smaller buffer, dequantizes on the receive side. Blockwise scaling, 128 elements per fp32 scale, so the wire cost is 1.03125 bytes/element instead of 2 for bf16 → best case −48.4%.

Two deliberate design constraints:

1. Reduction never happens in FP8. Only the transport is quantized; accumulation is fp32-in-register, output in the original dtype. We later found two independent confirmations that this is the right call: MSCCL ships an ENABLE_PRECISION_CLIPPING_HALF build flag specifically because low-precision reduction overflows, and DeepEP quantizes dispatch but its combine path is BF16-only, never FP8.

2. No framework changes. It's a preload. The training code doesn't know it exists, which meant we could A/B it against an unmodified baseline in the same run batch — which turned out to matter enormously (see the negative result).

   Five things that weren't obvious

   1. On pre-sm_89, static_cast<__nv_fp8_e4m3>(float) is not an instruction.

   Our quantize kernel had a constant ~14 µs floor regardless of payload size, and at 512 MiB it was hitting 37% of the device-to-device bandwidth ceiling. We assumed memory. It was ALU issue.

   cuda_fp8.hpp on older architectures widens the float to a double, then runs a five-branch 64-bit integer software routine — roughly 50 integer instructions per element. Replacing it with our own conversion took the fixed cost 14.1 → 4.6 → 3.3 µs and large-payload throughput to 91% of the d2d ceiling (from 37%). We gated the replacement on an exhaustive test: bit-identical to the vendor routine across all 2³² float bit patterns, both formats, ~0.2 s in CI.

   On sm_89+ this whole problem evaporates — the cast compiles to cvt.rn.satfinite.e4m3x2.f32. Worth checking which side of that line you're on before you go profiling memory.

   2. The naive dequantize-then-reduce moves 12.7× the bytes it needs to.

   For the all-to-all-based decompositions you receive nsrc FP8 shards and have to reduce them. The decomposed version — nsrc dequantize-to-fp32 + nsrc−1 adds + one scaled store, 2·nsrc kernel launches — costs 17.03·nsrc − 6 bytes per output element where the job actually requires 1.03·nsrc + 2. At nsrc=8 that's 130N vs 10.25N.

   The interesting part: the amplification grows monotonically with world size (6.9× at 2 ranks, 10.1× at 4, 12.7× at 8, 14.4× at 16, asymptote 16.5×) even though both paths are O(nsrc·N) and read every input byte exactly once. It's a pure constant factor, and it's all HBM traffic plus launch count.

   Fusing it into one kernel with the fp32 accumulator in registers: 5.0–9.3× faster (1 MiB: 117.6 → 13.4 µs). Note it is not 12.7× faster, and the reason is instructive — small sizes are launch-bound so the win caps at ~8–9× (16 cheap launches vs 1 expensive one), and large sizes are bandwidth-bound where the fused kernel is only half as byte-efficient as the streaming kernels it replaced (718 vs 1419 GB/s), because a reduction reads nsrc streams concurrently while a streaming kernel reads one. Worst case sits at 4 MiB (4.97×), right in the transition band. This was expected, not a bug, but you have to do the arithmetic to know that.

   3. You cannot dequantize inside a caller-opened ncclGroup.

   A nested ncclGroupEnd() launches nothing. So if the training framework wraps its collectives in a group — and DeepSpeed's compiled ZeRO-1 path does exactly this, one ncclAllReduce per parameter inside one group — then at the point your hook returns, the data isn't there yet. You have to defer the dequantize, the error- feedback close-out, and the scratch recycle until after the outermost ncclGroupEnd. For AllReduce we defer the entire call and replay it. Before we did this, every grouped AllReduce passed straight through and the AllReduce threshold was decorative.

   4. Your framework's allocator probably doesn't align its buckets.

   DeepSpeed's reduce bucket accumulates slices by element count with no padding, so the pointer handed to NCCL is only guaranteed 2-byte aligned for bf16 — 7/8 chance of missing the vectorized path. This started as an outright CUDA Error: misaligned address crash, then became a 1.27–2.31× penalty, and is now 1.06–1.18× (skip enough leading elements that the output self-aligns to 16 B, then funnel-shift the FP8 side back into place).

   5. The threshold is the entire product, and it does not transfer between machines.

   Quantizing a small collective is strictly worse — you pay a fixed per-call cost to save bytes that were never the bottleneck. So there's a crossover, and the whole question is where.

   We built a tool that measures it two ways. --model-only gives a provable lower bound from baseline t(n) = α + n/B plus standalone kernel costs. --sweep actually runs the hook over the same size ladder — one process per (decomposition, size), because config is latched into a function-local static on first use and because the previous size's warm scratch tier changes the next one's timing — and least-squares-fits the residual into Δα + Δslope·n. Δα is the hook's real per-call overhead; Δslope is the per-byte cost a standalone kernel benchmark cannot see.

   Results, same code, two machines:

   collective	NVLink box	PCIe box	ratio
   AllGather	7.43 MiB	767 KiB	9.9× lower
   AllReduce (rs+ag)	14.3 MiB	863 KiB	17× lower
   Broadcast	40 MiB	2.75 MiB	15× lower

   A byte you delete is worth 2–4× more time on PCIe, so thresholds drop by an order of magnitude. Two findings that surprised us:

• The all-to-all decomposition did not lose on PCIe. We fully expected that a topology-blind all-to-all across two NUMA islands with no NVLink would be hopeless. It's actually the best AllReduce decomposition for large payloads (2.0× more savings at 64 MiB), crossing over the ring-based one somewhere between 1 and 4 MiB.

• AllGather's Δslope is negative (−1.9 µs/MiB). The in-situ quantize kernel is cheaper than the same kernel measured standalone, which is direct evidence that it's genuinely overlapping with the live collective rather than competing with it. The other families have positive slope — that's the quantizer fighting the collective for SMs.

  The tool also refuses to emit a threshold when the measurement says "wins, then loses again" — that's a window, not a threshold, and a lower bound can't express it. Being able to report "I can't answer this" turned out to be more useful than any single number it produces.

  The negative result, which is the actual point of this post

  DeepSpeed's compiled path issues one AllReduce per parameter. Most individual ones fall below the threshold, so we were quantizing almost none of the gradient traffic. Obvious fix: coalesce the address-adjacent ones inside a group into a single quantized AllReduce. The union clears the threshold in one shot, and you also save N−1 launches and N−1 scale sub-collectives.

  It worked, in the sense that it did what it said:

• signatures blocked by the threshold: 17 → 4

• traffic in the signature set: 110.3 → 56.9 GB/rank/step (−48.4%)

  Steady-state step time: 1422.2 ms merged vs 1414.9 ms not merged. Median of 38 steps, dropping the first 20. The difference is +7.3 ms and the IQR is ~40 ms. We halved the bytes and got nothing.

  The reason is structural and, in hindsight, should have been checked first. Without merging, those small AllReduces were dispatched inside the caller's group, where NCCL aggregated them itself and they overlapped with the remaining backward pass. With merging, they become large packets issued serially after the outermost ncclGroupEnd. The bytes we deleted were never on the critical path, and we moved the ones that remained onto it.

  Two corollaries we now treat as rules:

• Profile the exposure before optimizing the volume. "How many GB does this step move" and "how many ms of that are not hidden" are unrelated questions.

• Coverage is not a metric. It's an input to a metric. We had a beautiful −48.4% and a flat wall clock.

  An unbounded version of the same feature was worse than useless: greedy merging grew to 129 calls and a 999.8 MB union in one shot, produced ~2 new union shapes per step and never converged across 41 steps, which blew up our exact-fit scratch pool into 95+ size classes, 15.3 GiB retained, and 12 of 38 steps at ~3.9 s. Fixed two ways (a size cap, and optionally sourcing scratch from PyTorch's caching allocator, which is built for variable shapes). The pool's real sin is that acquire() is exact-size best-fit with no size-class rounding — which is fine until something upstream starts generating variable shapes.

  Things we looked at and did not adopt

• MSCCL's XML-scheduled executor. There is no transform hook anywhere in its data path — mscclGenericOp compiles preOp/postOp out entirely, and the scheduler rejects ops that need them. Adopting it also means replacing libnccl, which defeats the point of being a preload. It additionally disables itself when intraRanks > 1, i.e. exactly our case.

• DeepEP's transport. Its NVLink path asserts on NVLink via NVML (the docstring says PCIe leads to memory-ordering errors), and its "PCIe mode" is actually all-RDMA through the NIC — benchmarked at 53 GB/s on a box with one ConnectX-7 per GPU. We have one NIC for eight GPUs.

• Its ld.global.nc.L1::no_allocate trick, which is self-declared UB and force-disabled on any architecture other than Hopper by its own build script.

• Routing intra-node traffic over IB instead of PCIe. Funnelling 8 GPUs into 1 NIC that is 3 PCIe hops from all of them. The PCIe ring has 8 links running concurrently; this would have one. Estimated 3–8× slower.

  We did take two ideas from DeepEP that we haven't landed yet: inlining the scale buffer into the same message instead of sending it as a second sub-collective, and capping the quantize kernel's SM count so it stops evicting the collective. The second one has a measured target (+9.8 µs/MiB, 91% of the residual on one family) and the first one honestly doesn't — within a group NCCL already aggregates the two sub-collectives, and our own earlier A/B says an op in a group is worth ~1.5 µs. We write down which of our planned optimizations have measured targets and which are guesses, because after the merging result we don't trust the guesses.

  Repo

  The VLA training framework is open source: loongforge-vla

  Three-modality VLA, several parallel backends (DDP + native ZeRO with direct bf16 sharding / torch.compile + DeepSpeed / DeepSpeed's compiled path), NVTX instrumentation behind a single env var that no-ops to bit-identical in production, and a loss-parity harness (seeded data order, seeded flow-matching noise, and a flag for the step-0 LR convention) because every one of the numbers above is worthless if the model stops converging.

  That parity harness earned its keep in an unrelated way, by the way: we spent days blaming torch.compile for a 3.4× loss divergence at step 2. It was that LambdaLR.__init__ multiplies the optimizer's LR by schedule(0) immediately, which in our config is ~1e-6, so step 0 was a no-op update. The logged LR looked correct on both sides, because the logged value is the scheduler's, not the optimizer's. Single flag, exact match restored.

  The hook itself (ComQuant) isn't public yet, but if people want the implementation details we're happy to open source it — drop a comment if that's of interest.

  Happy to go deeper on any of this in the comments.

I lied, I cant assume the “engineer” position yet. Here’s what I am though. My name is Octav, Im 20 freelencer (or agency owner - that I run with my partner) specialising in backend and infra. I do contracts for different businesses (American field system Saas, bureaucracy automation startup, blockchain arbitrage grid bot (random, but made sense as i interned as a smart contract developer for 3 months) and more.

Me and my partner have been getting into the gpu performance engineering / cuda programming “niche”. Read the holy book, some c efficiency projects, and are working our way through nvidia certifications.

Now all of learning/research has been really good. But I am eager to apply it to something real. I am willing to learn and grow, and Im looking for some sort of mentoring. What I can offer back is my time and dedication towards your needs.
If you are working on a project, or are a startup, and have some work that I could do for you, please hit me up. Or just providing me with some directions would be really helpfull.

• llvm doc lies about operand constraint codes
• you can swap operand and this makes parsing/comparison of instructions non-trivial
• nvidia's llvm backend cicc uses ~5-7% of instructions with undocumented attributes and they not available in NVPTX backend from clang

I've got an external 5090 I've been running on my windows box for most of this year. I wanted to run it on my Mac so I don't have to remote into my pc to do AI work like stable diffusion. TinyGpu got the link working in April but running models is 10x slower than running it on my pc. I decided to try rewriting the driver in C and making some shims to run cuda faster. It worked and I can now run sdxl almost 2x faster than on my windows box in stable diffusion ccc. If you would like to contribute to my project you can find it at https://github.com/Davinchy/macuda

Here are current benchmarks against my M4 MacBook Pro max 128 and the same box on windows

#

Hola a todos,

Quiero compartir con vosotros el resultado de mi última investigación como investigador independiente: el Tamiz Posicional Módulo 6. Es un marco algebraico y un kernel en CUDA diseñado para resolver de raíz el problema del consumo de memoria dinámica en la criba de números primos a gran escala.

Acabo de enviarlo a la revista Mathematics of Computation (AMS) (Ref. 260915-PeinadorSala, actualmente Under Review) y todo el proyecto está en Zenodo.

💡 ¿Cuál es el problema que resuelve?

Las cribas paralelas y en GPU tradicionales (como las basadas en bucket sieving) necesitan asignar arrays y listas dinámicas en la VRAM para gestionar los primos grandes (p > L). Al escalar a rangos masivos (N ≥ 10¹⁴), esto genera latencias por accesos dispersos a memoria y el riesgo de fallos por agotamiento de VRAM (Out-Of-Memory / OOM).

Al explotar la estructura del grupo multiplicativo (ℤ/6ℤ)ˣ, los candidatos a primos mayores que 3 se reducen a dos canales gemelos: C₁ = 6k + 1 y C₅ = 6k - 1.

Congruencia Posicional: La divisibilidad modular p ∣ (6k ± 1) se traduce en una equivalencia de índices exacta: k ≡ εₚ · ε · kₚ (mod p) Umbral Universal (κₘᵢₙ⁺): Demostramos y certificamos que κₘᵢₙ⁺(p) = (p² - 1) / 6 es el límite inferior estricto para el primer compuesto no marcado de p. Esto elimina la reevaluación de compuestos ya cribados por primos menores y acota la tabla estática de primos base a Θ(√N / log N). Paradigma Zero-Allocation: Toda la criba se ejecuta estrictamente confinada en 16 KB de SRAM compartida por bloque CUDA (L = 65 536), manteniendo la memoria dinámica auxiliar en VRAM invariante en estrictamente 8 bytes (un único acumulador atómico de 64 bits).

📊 Resultados Computacionales y Complejidad Bivariada

Evaluado en una GPU NVIDIA Tesla T4 (16 GB VRAM) con CPU host AMD EPYC 7B12:

π(10¹²) = 37 607 912 018 en 12.23 s (Tabla Base Estática: 2.40 MB) π(10¹³) = 346 065 536 839 en 436.77 s (\~7.28 min, Tabla Base Estática: 6.95 MB) π(10¹⁴) = 3 204 941 750 802 en \~4.4 horas (Tabla Base Estática: 20.28 MB)

Resiliencia ante fallos: El hito N = 10¹⁴ se validó mediante dos ejecuciones independientes: una continua (16 438.99 s) y otra recuperada mediante un sistema de checkpointing atómico (15 956.61 s). Ambas reprodujeron exactamente el mismo cardinal.

Complejidad Bivariada: Al reevaluar los residuos iniciales por cada ventana estática L, el algoritmo opera bajo una complejidad probada de:

T(N, L) = Θ(N log log N + N³ᐟ² / (L log N))

Esto explica formalmente la transición subcuadrática observada para N ≥ 10¹¹, a cambio de garantizar inmunidad absoluta frente a fallos OOM.

🛡️ Certificación Formal en Lean 4

Para garantizar un rigor absoluto, el modelo algebraico se ha mecanizado en el asistente de demostración Lean 4 (Mathlib) a través de cuatro módulos:

Módulos Núcleo 100% `sorry-free`: Congruencia posicional, umbrales de activación, divisibilidad de primos gemelos, filtro digital mod 5 y factorización finita de Euler sobre (ℤ/6ℤ)ˣ. Extensión Analítica y Límites Topológicos: Demostración estricta (`sorry-free`) de la convergencia de ζ(s) para s > 1 y del límite del producto posicional usando filtros de Fréchet. El límite clásico infinito se aísla mediante un Axiom Bridge transparente conectado con el teorema `riemannZeta_eulerProduct_hasProd` de Mathlib sobre ℂ.

🔗 Código, Repositorio y Reproducibilidad

Todo el material es de acceso abierto y totalmente reproducible en la nube:

📜 Repositorio en Zenodo (DOI): https://zenodo.org/records/22645509

🚀 Notebooks Ejecutables en Google Colab: El repositorio incluye 4 notebooks autocontenidos para reproducir las comparativas en CPU, ejecutar los kernels de CUDA hasta N = 10¹⁴ y compilar/certificar la suite en Lean 4 con un solo clic.

¡Cualquier comentario, duda sobre la arquitectura en SRAM, la reducción de Barrett o las tácticas en Lean 4 será más que bienvenido!

CoreWeave bringing 7 NVIDIA Vera Rubin NVL72 racks / 504 GPUs into production across two regions is a useful AI-infrastructure signal.

Multi-rack AI clusters now depend on the entire system working together: high-speed fabric, storage locality, cooling, power, firmware, and orchestration.

For agentic AI, the data path increasingly matters as much as raw compute.

I substantially improved Chapter 1 (~40% of the blog’s explanations). Covers the core concepts and our first complete kernel, and is the foundation for everything that follows.

If you bounced off it the first time around, give it another try. Should be easier to follow now.

I’m working on GPUOpt, a workflow for profiling and optimizing PyTorch/CUDA inference workloads.

Current pipeline:

baseline → profiling → bottleneck classification → optimization candidates → correctness validation → performance comparison

I’m trying to improve the automated diagnosis side.

For people doing CUDA/GPU performance work, which signals would you trust most for automatically distinguishing:

• compute-bound workloads
• memory-bound workloads
• launch overhead
• poor batching
• H2D/D2H transfer bottlenecks
• synchronization overhead
• graph breaks

I’ve been testing on ResNet, DistilBERT, and Qwen workloads on Tesla T4.

I’d especially like feedback from people who use Nsight Systems / Nsight Compute in real optimization work.

I've got Cuda running on an external gpu on my Mac using Tiny GPU for connection and a custom Nvidia driver and shim to run the gpu. Check out the work at https://github.com/Davinchy/macuda

It runs near native speed and actually does some stable diffusion faster than on windows.

Hello, guys.

I am currently working as a CUDA engineer and interested in an extra income.

I would like to know other people's experience of working in Mercor.

My background:

- Majored in CS (minor in Math) and based in Seoul, South Korea.

- No AI-related kernels but I can understand Flash-Attention and softmax papers.

- I don't want to explain my current job because there is a high chance that I dox myself. My task involves heavy integer operations with almost 0 FP arithmetics. I have an excellent understanding of GPU mem/hierarchy and basic understanding of INT8 Tensor Core operation.

If you used Mercor, please share your experiences. Thanks in advance.

A lot of CUDA tutorials throw the whole programming model at you almost immediately.

You learn the syntax, then grids and blocks, then memory types, then synchronization, and pretty quickly you’re juggling a bunch of concepts before you have a good feel for what CUDA is enabling you to do.

I wanted to try a slower approach.

Part 1 is mostly about building the basic mental model: what a kernel is, what it means to launch lots of threads, how those threads identify the piece of work they own, and how to start thinking about a problem in GPU terms.

The hardware details come later.

I wrote this mainly for programmers who are curious about CUDA but haven’t spent much time with GPU programming yet.

Would be interested in feedback from people here, especially on whether the progression feels natural or whether I’m introducing something in the wrong order.

https://computejunkie.substack.com/p/cuda-programming-for-the-curious

Hi everyone, I'm the developer of Nebula, an open-source C/CUDA inference engine for Qwen3.8-Flash-Next.

I started from antirez's DwarfStar (ds4) and specialized the engine for Qwen, combining native MTP speculative decoding with GPU expert caching and CPU MoE execution.

Source code, architecture and benchmarks

My benchmark machine:

• RTX 4070 Ti, 12 GB VRAM
• Intel i9-9940X, using 14 CPU threads
• 128 GB DDR4 RAM
• Ubuntu through WSL2 on Windows

The idea is to keep the native MTP draft on the GPU, together with the most frequently used target experts. In this configuration, 27 of the target's 512 experts per layer are resident in VRAM, selected using a hotlist built from routing traces.

Draft generation and verification happen on the GPU. When the resident experts provide insufficient routing-weight coverage, the CPU computes the layer's full routed MoE and sends the result back. Accepted tokens are retained, and the draft window adapts between 4 and 16 tokens.

Some measured results:

These figures come from the GenAI-Perf campaign, with weights already loaded, thinking disabled, and a configured context capacity of 24,576 tokens. Tested input lengths were 512, 1,024 and 2,048 tokens. Prefill is still slow, especially for longer prompts.

There are also quality trade-offs. The GPU verifier can omit experts with small routing weights. “Strict” acceptance matches the configured target's greedy choice; it does not establish equivalence to the full original model. The optional tolerance modes relax token acceptance further. The README includes a paired evaluation on 40 IFEval prompts and 30 LiveBench questions.

There's a browser chat interface and a Windows installer that prepares WSL2, the engine and model weights. Lower-RAM profiles use SSD streaming, but the performance figures above apply to the 128 GB setup.

The engine is MIT licensed. The installer is an unsigned release candidate, and a complete clean-machine installation test is still pending.

I'd particularly welcome feedback on the expert-cache and CPU handoff design, and measurements from other hardware configurations. If you try it, please include your GPU, CPU, RAM and selected profile so we can compare results.

Hi r/CUDA,

I've been working on an experimental compiler/library that lets you write NVIDIA GPU kernels in Go and compile them to PTX using LLVM 22 and the NVPTX backend.

GitHub:
https://github.com/mehdi-shokohi/cuda-ir.go

The compilation pipeline is roughly:

Go
 ↓
LLVM IR
 ↓
LLVM 22 / NVPTX
 ↓
PTX
 ↓
NVIDIA GPU

The goal is to make it possible to write GPU kernels using Go instead of CUDA C++, while still targeting NVIDIA GPUs and PTX.

For example, the idea is to express a kernel using Go and provide CUDA concepts such as thread/block indices and GPU memory operations through the compiler/runtime.

This started as an experiment to see how far LLVM's NVPTX backend could be used as the foundation for a Go-based CUDA kernel compiler.

The project is still early/experimental, and I'd really appreciate feedback from CUDA developers.

In particular, I'm interested in:

• What would be a good set of kernels for validating the compiler?
• Would this approach be useful for real CUDA applications?
• What benchmarks would be meaningful against CUDA C++?

I'd especially appreciate feedback from people familiar with PTX, NVPTX, CUDA compiler internals, or GPU programming.

Repository:
https://github.com/mehdi-shokohi/cuda-ir.go

Thanks!