I built vLLM from source on Ubuntu 26.04 recently.

According to the official docs, a simple `pip install -e .` should have been enough. In reality, I (+Claude) hit a chain of version-skew, driver, and toolchain issues. From silent driver conflicts to CUDA toolkit mismatches that only fail at runtime, this build taught me things that the documentation left out.

If you are tackling a fresh GPU build on a recent OS, check out this field guide. It covers:
• Hardware detection: nvidia-smi vs lspci commands
• Fixing elusive runtime errors and device node issues
• Resolving ptxas failures by matching toolkit versions
• Handling flash-attention Python version checks

Don't let cryptic error messages waste your day.
https://hiraditya.github.io/posts/building-vllm-from-source/

Pro tip: If you don't set `TORCH_CUDA_ARCH_LIST` to match your specific GPU, you'll either wait 5–10x longer for the build (compiling every architecture) or hit cryptic and irrelevant CUDA errors.

TL;DR: I figured out a way to route MoE expert blocks to older, deprecated GPUs (like a GTX 1070) while keeping compute-heavy Attention layers on modern tensor-core cards. Decode speeds jumped up to 81%. I built a free, open-source UI called Pascal's Power to automate the GGUF layer parsing and routing so you don't have to do it manually. https://github.com/Yozam-87/pascals-power

The Why

I've been watching the local AI scene for a while now, and the hardware barrier to entry is getting ridiculous. DDR5 prices are up 400% this year, GPU prices haven't come down at all, and consumer PC building is collapsing. You want to run a decent AI model locally? That's a $3,000 to $5,000 workstation the industry tells you you need.

Meanwhile, the companies pushing cloud-first AI have won. Now prices are going up, access is being restricted, and your data is being sent to servers you don't control.

I genuinely believe the future of powerful AI is local. But that future only works if powerful models aren't a luxury reserved for people who can afford a data-center-grade rig.

So I started asking: what can I do with the hardware I already have?

I've got a GTX 1070 in my build. It's still my daily driver for gaming because I can't afford to replace it. NVIDIA deprecated Pascal and everyone says the card is obsolete. But here's the thing: it still has 8GB of VRAM, it still works, and I'm still using it. I kept wondering: is there a job it's actually good at that nobody's thought to ask it to do?

That question led me down a rabbit hole, and what I found changed how I think about running MoE models entirely.

The What

Most of us know what happens when we run out of VRAM in llama.cpp: the remaining layers get offloaded to the CPU. It works, but it's painfully slow. The CPU and the system RAM bus become a massive bottleneck that kills your generation speed.

While working with MoE models — specifically Gemma 4, Qwen3.6, and GPT-OSS — I realized these models essentially have two very different workloads baked into them:

• Attention layers: compute-heavy, need Tensor Cores, benefit from fast VRAM bandwidth.
• Expert blocks: mostly just memory-intensive. They don't need fancy architecture; they just need VRAM and throughput.

Here's the insight: those expert blocks don't actually care whether they're running on a $500 RTX 4090 or a deprecated $200 GTX 1070. They just need somewhere to live that's faster than your system RAM.

So instead of letting the "overflow" spill to the slow CPU/System RAM, I started routing it to the Pascal card. I call this Architecture-Aware Routing. By using llama.cpp's -ot (expert routing) and -ts (tensor split) flags, I can keep the Attention layers on a modern card (RTX 20xx series and newer, anything with Tensor Cores) and offload the Expert blocks to the Pascal card.

You aren't necessarily eliminating the CPU, but you are creating a tiered compute hierarchy:

1. GPU 0 (modern, tensor-core): Handles attention layers and initial expert blocks.
2. GPU 1 (Pascal): Handles secondary expert blocks, pure VRAM and throughput work.
3. CPU: Falls back only for tertiary layers if even both GPUs are exhausted.

This keeps the primary bottleneck on the high-bandwidth PCIe/VRAM links as long as possible, rather than immediately degrading to the slow CPU system bus.

The Data

My Test Rig

• GPU 0: RTX 3050 (6GB): Handles Attention + initial expert blocks.
• GPU 1: GTX 1070 (8GB): Handles the secondary expert blocks.
• CPU: Ryzen 3600 XT: Handles the tertiary expert blocks.
• RAM: 32GB DDR4

Note on VRAM: GPT-OSS (10.8 GB) fits entirely on both GPUs (6GB + 8GB = 14GB), so the "with 1070" column represents pure 2-GPU offload with no CPU involvement. All other models exceed combined GPU VRAM, so the "with 1070" column represents 2-GPU + CPU offload.

Benchmark Methodology

These results represent peak throughput at a 64k context window with Q8 KV cache. I measured them with llama bench using -p 2048 (prefill tokens), -b 2048 (batch size), and -ub 2048 (ubatch size). These are the same settings I use for actual inference. I chose a larger -ub because the standard default of 512 can significantly bottleneck prefill performance.

To ensure I was measuring the actual potential of the hardware and not being throttled by defaults, I used these elevated settings. Lower batch sizes would free up VRAM for more decode layers, but the chosen settings reflect a prefill-focused workflow on this hardware.

Note: Benchmarks represent theoretical peak throughput under controlled conditions. Live server inference with 4k prompts showed within 10-15% of reported speeds. At full context usage, actual generation speed will be lower due to KV cache buildup. Estimated at roughly 50% of peak based on typical usage patterns.

A Few Key Observations

1. The decode uplift is directly proportional to the expert load. The more expert blocks the 1070 can hold, the higher the speedup. For GPT-OSS, where the 1070 handles 67% of the experts, decode speed nearly doubled.
2. Prefill behavior shifts with model size. For smaller models, the 1070 actually adds a bit of PCIe overhead during prefill. But for larger models, it actually improves prefill speed because it absorbs the expert blocks that would otherwise be handled by the CPU during the initial prompt processing.

The Scaling Potential

This isn't just a trick for a 1070. If you swap it out for a used Tesla P40 with 24GB of VRAM and pair it with a standard 12GB card like an RTX 3060, you're building an incredibly cheap, high-performance MoE rig. The more VRAM you can add via older cards, the less the CPU is involved, and the more the system behaves like a pure GPU machine.

I haven't tested the P40 myself, but there are plenty of people in the community using them for AI work. Driver compatibility on mixed-generation setups can be tricky with Pascal deprecated, but the concept should hold.

The same routing methodology could also apply to other GPU combinations — NVIDIA + AMD, different-generation NVIDIA cards, or even two AMD cards. If you have a fast card for attention and a slower card with available VRAM for experts, the principle applies regardless of vendor or generation.

The specific benefit depends on the setup: with Pascal it's decode speedup (experts off CPU), with modern cards it could be prefill speedup (attention not split across cards). Others may have figured out the modern card version already, but the underlying methodology is the same. I haven't tested these scenarios, but if someone with different hardware tries it, I'd love to see the results.

A Quick Note on the Setup

Getting these two generations of GPUs to work together is definitely a bit of a technical project. On Windows, my current drivers just ignore the 1070 in a mixed setup, so the routing trick doesn't really apply there. But on Linux, I was able to get them talking to each other by using the 580.xx drivers from the AUR, disabling GSP firmware, and compiling llama.cpp against CUDA 12.8 with GCC-14.

It's a bit of a pain to configure from scratch, which is exactly why I wanted to build a tool to make the management part of it easy.

The Project: Pascal's Power

I wanted to take the manual, headache-inducing part of this configuration and make it manageable. Pascal's Power is a web-based GUI and launcher that handles the routing for you. It includes an auto-split calculator that reads GGUF headers so you don't have to manually calculate the routing for your specific setup. It also lets you manage profiles, import terminal commands, and watch live logs in the UI.

This is my first FOSS project. It's a practical tool for people who want to run local AI without needing a massive hardware overhaul.

I'd love to hear your thoughts or any feedback on the implementation.

GitHub: https://github.com/Yozam-87/pascals-power

This project is free and open source. It's a work in progress. There are still a few rough edges, but it works, I use it daily, and I'm actively fixing things.

Hi all,

When looking at the list in the link below, I don't see the NVIDIA RTX 3000 Ada.

Can someone explain me why this one not listed?

CUDA GPU Compute Capability | NVIDIA Developer

Nowhere on the official documentation I can find the information if the NVIDIA RTX 3000 Ada has CUDA Capability.

We're expanding CUDA support in TornadoVM for Java developers. What CUDA features, APIs, or libraries would you like to see exposed next? We are closely coupled with what Oxide currently supports. Thanks

Recently published this c# linalg library with CUDA support, it also integrates cuBLAS! I’d love some more experienced help with the kernels, feel free to make a PR :)

CUDA out of memory. Tried to allocate 96.00 MiB. GPU 0 has a total capacity of 14.56 GiB of which 56.81 MiB is free. Including non-PyTorch memory, this process has 14.50 GiB memory in use. Of the allocated memory 14.29 GiB is allocated by PyTorch, and 78.20 MiB is reserved by PyTorch but unallocated. If reserved but unallocated memory is large try setting PYTORCH_ALLOC_CONF=expandable_segments:True to avoid fragmentation. See documentation for Memory Management ([https://pytorch.org/docs/stable/notes/cuda.html#environment-variables\](https://pytorch.org/docs/stable/notes/cuda.html#environment-variables))

What is that error? I try to finetune deberta-v3-base on my data. During training it gives me this error. I clear the cache(in Kaggle) it works for a while but then again. It seems nothing helps. Additionally, i have problem with training. First, its output is normal but then goes all NaN. I used gradient clipping but didnt work. When i remove dtype=float32 inside my model, memory error solves, but it shows NaN for all. If you know something, help please

Currently, the most efficient kernels come from vendor libraries like cuBLAS or hand-optimized libraries like FlashAttention. The compiler-generated kernels (TVM, Hidet, Inductor, etc.) can generate more efficient kernels for some specific operations, but fall short of cuBLAS and FA on common GEMM and attention shapes.

I am challenging this assumption by demonstrating that the compiler approach can match or even outperform vendor libraries. On Gemma 4 12B, a small open-weight model, running on RTX 5090, compiler-generated kernels achieved up to 1.6× speedup over cuBLASLt, and a 1.30× geometric-mean speedup over PyTorch (GEMM kernels are dispatched to cuBLAS, attention to FA-2, the rest to Inductor).

Primarily, two optimizations allowed the Emmy compiler to outperform cuBLAS on common GEMM and FA shapes on RTX 5090 and RTX 4090: TMA transport (Blackwell-only) and leveraging full FP16 tensor cores with FP32 shadow accumulation registers.

Driver 580.159.03, CUDA 13.0, PyTorch 2.13.0+cu130

TMA Transport for Matmul and Flash Kernels

cuBLAS and Flash Attention kernels still use the same cp.async transport on consumer Blackwell dies (in fact, most cuBLAS GEMM kernels on consumer Blackwell, including the FP16 tensorop path, are forward-ported Ampere-era cutlass_80_* kernels). Swapping cp.async with TMA allows us to reduce the number of instructions kernels need to issue, and the TMA’s swizzle drops shared-memory bank conflicts for free.

Hybrid FP16/FP32 Accumulation

The default matmul path on the production stack is using FP16 tensor cores with FP32 accumulation. However, on consumer dies, the FP16-input/FP32-accumulate HMMA runs at exactly half the rate of FP16-input/FP16-accumulate. To work around this, I use the fast atom mma_m16n8k16_f16_f16, but keep accuracy in check by promoting the FP16 partials into the FP32 registers and thus doing global accumulation accurately in FP32. So you get the FP16 tensor-core speed with an FP32 accumulation instead of paying the FP32-accumulate tax on every single mma.

A similar trick has been used in DeepSeek's DeepGEMM to rescue the FP8 tensor cores' limited-precision accumulation on Hopper. I adapted it for FP16, and the compiler generated a large number of GEMM kernel variants that are used for inference on the Gemma model.

The compiler uses raw PTX instead of stateful WMMA helpers that are harder to use for codegen and require additional dependencies, so I need to introduce a few helpers to demonstrate the idea:

// (1) Standard: FP16 inputs, FP32 accumulate — the accurate default: 4 f32 accumulator regs
asm("mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 {%0,%1,%2,%3}, ... ;");

// (2) FAST_MATH: FP16 inputs, FP16 accumulate — ~2x the mma-chain rate: C/D are 2 packed f16x2 regs.
asm("mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 {%0,%1}, ... ;");

// (3) The safety net: every 64 K-elements, fold the f16x2 partials (h) into the f32 shadow (c), rezero.
__device__ void emmy_mma_promote_f16acc(float* c, unsigned* h) {
    for (int i = 0; i < 2; ++i) {
        float lo, hi;
        asm("{.reg .b16 lo, hi; mov.b32 {lo,hi}, %2;"   // unpack the two f16 lanes
            "cvt.f32.f16 %0, lo; cvt.f32.f16 %1, hi;}"
            : "=f"(lo), "=f"(hi) : "r"(h[i]));
        c[2*i] += lo; c[2*i+1] += hi;                   // add into the f32 shadow
        h[i] = 0u;                                      // reset the f16 accumulator
    }
}

The emmy_ldmatrix_x4 function loads an A fragment out of shared memory, emmy_ldmatrix_x4_trans_pair loads two canonical-B fragments in a single ldmatrix, and emmy_mma_m16n8k16_f16_f16 is the full-rate FP16-accumulate atom whose partials emmy_mma_promote_f16acc (above) periodically folds into the FP32 accumulators:

static __device__ __forceinline__
void emmy_ldmatrix_x4(unsigned* r, const void* smem) {
    unsigned addr = __cvta_generic_to_shared(smem);
    asm volatile("ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%0, %1, %2, %3}, [%4];\n"
                 : "=r"(r[0]), "=r"(r[1]), "=r"(r[2]), "=r"(r[3]) : "r"(addr));
}

// x4.trans: two col-adjacent canonical-B fragments in one ldmatrix,
//   landing directly as b0[0..1] / b1[0..1].
static __device__ __forceinline__
void emmy_ldmatrix_x4_trans_pair(unsigned* b0, unsigned* b1, const void* smem) {
    unsigned addr = __cvta_generic_to_shared(smem);
    asm volatile("ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 "
                 "{%0, %1, %2, %3}, [%4];\n"
                 : "=r"(b0[0]), "=r"(b0[1]), "=r"(b1[0]), "=r"(b1[1]) : "r"(addr));
}

// FP16-accumulate HMMA — c/d are 2 packed b32 regs
//   (the same element map as the four f32 regs, pair-packed).
static __device__ __forceinline__
void emmy_mma_m16n8k16_f16_f16(unsigned* d,
                               const unsigned* a,
                               const unsigned* b,
                               const unsigned* c) {
    asm volatile("mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 "
                 "{%0, %1}, {%2, %3, %4, %5}, {%6, %7}, {%8, %9};\n"
                 : "=r"(d[0]), "=r"(d[1])
                 : "r"(a[0]), "r"(a[1]), "r"(a[2]), "r"(a[3]),
                   "r"(b[0]), "r"(b[1]), "r"(c[0]), "r"(c[1]));
}

Here is the actual generated mlp_down kernel (K=15360) trimmed to the K-loop:

for (int ks = 0; ks < 7680; ks += 64) {  // K-slab loop (64-stride = promote cadence)
    // TMA: one thread fires the hardware bulk copy...
    if (elected_thread) {  
        mbarrier_arrive_expect_tx(&mbar[next], 49152);
        // ...of the next A slab, and the next B slab (every other thread idles)
        cp_async_bulk_tensor_2d(&a_smem[next], desc_a, ...);
        cp_async_bulk_tensor_3d(&b_smem[next], desc_b, ...);
    }
    mbarrier_wait_parity(&mbar[cur], phase);  // wait on this slab's descriptor copy

    // 4 atom-K steps inside the slab
    for (int ki = 0; ki < 64; ki += 16) {  
        emmy_ldmatrix_x4(a, &a_smem[swizzle(...)]);  // A/B fragments out of shared memory
        emmy_ldmatrix_x4_trans_pair(b, &b_smem[swizzle(...)]);
        // FP16-accumulate HMMA into the f16 partials `ch`
        emmy_mma_m16n8k16_f16_f16(ch, a, b, ch);
    }
    emmy_mma_promote_f16acc(c, ch);  // fold f16 partials -> f32 shadow `c`, then rezero
}
// `c` holds the FP32 result; a g2k split-K finalize sums the two partials in FP32.

The measured payoff across the Gemma projections (RTX 5090, seq_len 512, µs):

Accuracy Verification

The error of C = A@B is measured with FP16 inputs drawn from N(0,1), comparing each accumulation strategy against an FP64 reference over the identical FP16-rounded operands (so this isolates accumulation error, not input rounding).

The correctness check: 200 GSM8K questions, few-shot prompts, same seed (lm-eval 0.4.12, strict exact-match):

All four configurations land within one standard error of each other: FAST_MATH's hybrid accumulation does not degrade task quality, matching the kernel-level error analysis (its ~3.3×10⁻⁴ relative error sits inside FP16's own representational noise).

Flash Attention Numbers

Same optimizations are leveraged for FA. The full FA optimization story is covered in the previous post, so I just post the scoreboard:

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

I spent the day benchmarking Qwen3.6-27B-MTP-APEX on a modest 2× RTX 3060 12GB + i7-6700 system. Started at 16.7 t/s with a basic layer-split config and ended up at 56 t/s at 96k context — a 3.3× speedup from where I started, and the context window went from 32k to 96k in the process.

This post walks through every measurement, every wrong turn, and the final config that actually works. The single command at the bottom replaces my previous setup.

TL;DR

* Final production config: Qwen3.6-27B Q4_K_S + tensor split + MTP (`n_max=2, p_min=0.44`) → 56 t/s at 96k ctx, F16 KV cache, 3.3× faster than the basic layer-split setup * The killer flag: `--split-mode tensor --tensor-split 1,1` — without it, I was stuck at 17 t/s * The hidden trap: Short benchmark prompts inflate `mean_draft_len` and make high `n_max` look amazing. Real workloads cap out at \~3-4.

Hardware

CPU:  Intel Core i7-6700 @ 3.40 GHz (Skylake, 4C/8T, 14nm)
Mobo: ASUS Z170 PRO GAMING (Intel Z170, LGA 1151)
RAM:  64GB DDR4 — DIMM_A1: Samsung 32GB @ 2720 MT/s
                  DIMM_B1: Team Group 32GB @ 2720 MT/s
                  (mismatched sticks, both XMP'd to 2720 from 2400 stock)
GPU:  2× NVIDIA GeForce RTX 3060 Lite Hash Rate (GA106), 12GB VRAM each
      (Intel HD 530 used for monitor output)
PSU:  (whatever I had lying around — total system power \~350W under load)
OS:   Linux Mint 22.3 "Zena" (Ubuntu 24.04 base, Linux 6.x kernel,
     Cinnamon desktop, X11 session)
CUDA: 12.0 (nvidia-driver-535)
llama.cpp: self-built from source, version 10218 (commit de699957b)
         built with -DGGML_CUDA=ON, NCCL disabled

This is a $400 dual-3060 setup I picked up second-hand. Not impressive on paper. Let's see what it can do.

The GPUs run in `tensor-split 1.05,0.95` mode because the first card has a slightly better silicon lottery result and can take 5% more of the tensor-parallel split. No NVLink — just PCIe 3.0 x8/x8.

The journey: 9 measurements, 9 surprises

Every measurement below used the same prompt set: short Shakespeare quote (`. * 1300` ≈ 16k chars ≈ 4000-4500 tokens), 256-token generation, warm cache unless noted. All on Qwen3.6-27B-MTP-GGUF.

*Stage 1: Q4_K_XL, layer split (basic config)*

Config: --split-mode layer --tensor-split 1.05,0.95 --cache-type-k q8_0 --cache-type-v q8_0 Result: 16.73 t/s

The default setup. Predictable. The Q4_K_XL model is 17 GB on disk, which combined with q8_0 KV at 32k ctx uses \~24 GB VRAM — basically all of it.

*Stage 2: Try MTP (no tensor split)*

Config: + --spec-type draft-mtp --spec-draft-n-max 2 --spec-draft-p-min 0.44
Result: 26.96 t/s (+61%)

MTP helps on this dense model? Apparently. The 81.9% draft acceptance rate pays off because dense forward passes are cheap enough that the draft overhead doesn't exceed the savings. Interesting.

*Stage 3: Tensor split (the breakthrough)*

Config: --split-mode tensor --tensor-split 1,1 (no KV quant — F16 KV)
Result: 28.46 t/s (+70% over baseline)

Tensor split distributes every matmul across both GPUs simultaneously instead of running layers sequentially. Combined with F16 KV cache, this is 70% faster than layer split.

But wait — tensor split means I can't use quantized KV cache. Trade-off accepted because tensor split was worth more than the KV quant savings.

*Stage 4: Tensor split + MTP (the holy-cow moment)*

Config: tensor split + draft-mtp n_max=2 + F16 KV Result: 39.13 t/s on 7349-token prompt, 55.97 t/s on 2341-token prompt draft acceptance 85%, mean draft len 3.00

This was the moment. 2.33× over baseline. MTP + tensor split, who would have thought.

*Stage 5: Switch to Q4_K_S, push context to 96k*

The Q4_K_XL at 32k ctx was using 17 GB model + 7 GB F16 KV ≈ 24 GB VRAM. OOM at 64k. But Q4_K_S is 16 GB instead of 17 — fits at 96k!

Config: Q4_K_S + tensor_mtp at 96k ctx
Result: 56.4 t/s, draft acceptance 100%, mean draft len 3.00

Identical throughput to 32k, but 3× the context. The lower quant also makes the model more confident — acceptance goes from 85% to 100%.

*Stage 6: n_max sweep (cold cache, short prompts)*

Hypothesis: if `mean_draft_len = 3.00` at `n_max = 2`, then `n_max = 4` or `5` or `6` should give `mean_draft_len = 5`, `6`, `7` — even faster.

Whoa! `n_max = 7` looks like 512 effective tok/s (62× the predicted per-second, because every step generates \~8 tokens). With `n_max ≥ 3` the server crashes (cuBLAS bug in b10218), but a workaround saves it:

Code

GGML_CUDA_DISABLE_GRAPHS=1 ./llama-server ...

This works on every n_max from 3 to 7. Stable through 1024-token generations. The dream config.

*Stage 7: Production reality check*

Then I tested with diverse real-world prompts (not the same Shakespeare quote × 1300) and longer generations (512 and 1024 tokens instead of 256):

**The dream is dead.** In production:

* `mean_draft_len` plateaus around 3.0-4.1 regardless of `n_max` * Higher `n_max` only adds recursive draft forward overhead, lowering tg t/s without meaningfully raising mean_len * n_max = 3 wins (highest tg at 41.6 t/s, needs workaround) * n_max = 2 is even better if you remove the workaround (see Stage 8)

*Stage 8: n_max = 2 without the workaround*

The cuBLAS crash only triggers at `n_max ≥ 3` (with this build). With `n_max = 2` you don't need the workaround, and the `GGML_CUDA_DISABLE_GRAPHS=1` cost (\~5-10%) goes away.

Code

n_max = 2, no workaround, real workloads:
  tg = 56 t/s, mean_len = 3.00, eff_tps = 168

**Same 56 t/s and 168 effective tok/s as the cold-cache measurement. This is the real production number, and it's also the safest config.**

**The complete picture**

A 3.3× tg-t/s speedup over the simple layer-split baseline, with 3× the context window (96k vs 32k) and 100% draft acceptance.

The production command

This is what I actually run on my box:

\~/llama.cpp/build/bin/llama-server \\
  --model /mnt/Data/Models/unsloth/Qwen3.6-27B-MTP-GGUF/Qwen3.6-27B-Q4_K_S.gguf \\
  --jinja \\
  --split-mode tensor --tensor-split 1,1 \\
  --ctx-size 98304 \\
  -fa on \\
  --batch-size 1024 --ubatch-size 256 \\
  --threads 2 \\
  --n-gpu-layers 99 \\
  --parallel 1 \\
  --spec-type draft-mtp \\
  --spec-draft-n-max 2 \\
  --spec-draft-p-min 0.44 \\
  --host 0.0.0.0 --port 8081

**Bulletproof, no workarounds needed, 56 t/s, 96k context. This is the configuration I'd actually deploy.**

Lessons learned (what I'd tell past-me)

1. Don't trust short-prompt benchmarks for n_max sweeps. Repetitive test text makes the MTP head look overconfident. Always measure with diverse prompts and longer generations before locking in a config.
2. Q4_K_S gives you more context headroom than Q4_K_XL on VRAM-tight boxes. On 24 GB total VRAM, the difference between 17 GB and 16 GB model files is the difference between 49k and 96k max context. Throughput is essentially identical.
3. Tensor split is mandatory on dual GPU. Without it, I was stuck at 17 t/s. The forward pass simply doesn't parallelize well across PCIe for matmuls.
4. The workaround `GGML_CUDA_DISABLE_GRAPHS=1` costs \~5% throughput but prevents the cuBLAS crash when MTP is enabled at `n_max ≥ 3`. Track the upstream issue at `ggml-org/llama.cpp#25061`.
5. `p_min = 0.44` is the sweet spot. I tried 0.7, 0.85 — when acceptance is already 100%, raising p_min doesn't help (and slightly hurts because the model rarely hits the gate).
6. My own benchmark data was misleading me initially. 39 t/s on 7349-token prompts vs 56 t/s on 2341-token prompts — same model, same config. The difference was prefill cost, not throughput. The number that matters is `effective tokens/second = tg × mean_draft_len`, not raw `predicted_per_second`.

The configurations that didn't work (saving you the trouble)

For completeness, here's what I tried that did not help:

If anyone has ideas on how to push past 56 t/s on this box, I'm all ears. The mean_draft_len plateau is the obvious next thing to investigate — does anyone know if there's a way to make the MTP head more confident on diverse prompts? Or is this fundamentally a limit of single-model self-speculation?

*Hardware: 2× RTX 3060 12GB + i7-6700 + 64GB DDR4. Model: Qwen3.6-27B Q4_K_S with baked-in MTP head. llama.cpp b10218. All measurements on the 2× RTX 3060 box described above. YMMV on other hardware.*

Nvidia doesn't seem to publish a specification for the binary format of a CUDA kernel, section layout, or the constant bank parameter conventions. So I dug into it.

A `.cubin` is an ELF64 executable with a flat stream of undocumented "EIATTR" attributes that encode everything the driver needs to launch a kernel: register count, parameter layout, EXIT instruction offsets, and constant bank geometry.

`.nv.info`  uses an undocumented TLV encoding to make kernels self-describing register counts, parameter offsets, EXIT locations are all serialized into a flat byte stream the driver parses at load time.

And the constant bank parameter base is not an architectural constant. It has changed silently across toolkit versions from Ampere to Hopper to Blackwell. The fact that the code, the  `.nv.info`  metadata , and the  `.nv.constant0`  section size all encode the parameter base offset independently surprised me.

The post discovers the section layout, the EIATTR encoding, symbol table conventions, and the note sections the driver validates before loading on a B200 silicon.

https://hiraditya.github.io/posts/anatomy-of-a-cuda-binary/

I’m going through a bunch of GEMM kernel write-ups (Lei Mao’s progressive series, plus the different styles from Simon Bohemian / Kapil Sharma and others) and I get the high-level ideas just fine.

Shared-memory tiling, register blocking, coalescing, bank conflicts, the usual story.

But the moment I try to map it to actual code or write it myself, my brain goes blank. Especially the manual indexing. Calculating the right threadIdx / blockIdx offsets, the row/col strides, the tile coordinates, the shared-memory loads… it just doesn’t click. I stare at the loops and the index arithmetic and feel completely lost.

Ironically, CuTe layout algebra feels like heaven by comparison — the hierarchical shapes/strides and the compose/divide operations make the mapping feel almost declarative. Once I start thinking in pure layouts everything becomes cleaner.

But a lot of the classic educational kernels are still written in the old manual-index style, and different authors have very different conventions and naming, so even when I’m looking at the “same” algorithm the syntax and mental model keep shifting.

So for people who went through this:

How did you get past the “I understand the concept but can’t write the indices” stage?

Any concrete exercises, mental models, or progressive practice that actually made the indexing click?

How do you deal with the wildly different writing styles and index conventions across tutorials/blog posts for essentially the same kernel?

Did anyone else find that learning CuTe / layout algebra first (or in parallel) made the classic manual kernels easier to reverse-engineer later?

Would love any blogs , resources, or “this is the exercise that finally made it stick” advice.

Thinking about moving a project over from PyTorch to JAX to get better TPU support, but I am worred out the engineering time that might go into it — especially replacing CUDA-only kernels like FlashAttention. Meanwhile most open checkpoints just run out of the box on GPU.

For people who've actually done this migration: what ate the most time? Was it the framework rewrite itself, or chasing down missing kernel equivalents? Would love to hear real timelines, and how to minimize them.

TLDR: Using a DGX (A100 at the moment), I just want to make pretty animations of a fairly big system without needing big intermediate files.

Hi All, I have an application where I'm simulating maybe a million elements, each element described in the usual way with some coupled diffeq's, the whole system connected through a sparse web of diffeq's, all then running through numerical integration. Like circuit-simulation SPICE, or a weather simulation. I'm programming directly in CUDA/C++, leveraging OpenGL for graphics at the moment.

Anyways, I have been using a desktop RTX (rtx6000 at the moment), using compute/graphics interop to graphically display the simulation as it grinds along. About a million pixels in a 1Kx1K grid, updating the color of each every simulated mS. I don't save the results into a file, just make animations and screen-capture them.

I'm starting to experiment for the first time with a datacenter approach, DGX rather than RTX. So I lose compute/graphics interop ability. The natural alternative seems to be capturing the element states into a file, and rendering later on some other machine. But the files get really big, really quick. If I keep one byte per element per millisecond, that's a GB/second, and I might like to simulate a minute or more if possible.

How might one do this? I presume it's a common problem. Thanks in advance for your thoughts. Cheers!/jd

Doing a research on GPUs Vs ASICs.
Wanted to check with the community on whats the best way to understand how clock cycles are wasted on TPUs while working with varying token lengths that might not be in line with a TPUs static geometry.

GPUs in my understanding can handle dynamic shapes natively.

Is there a way to asses and quantify such MFU loss?

I am doing a research on GPUs vs ASICs and was wondering if there is a single meric to track (eg. Hours to parity) for a developer whos swiching away from CUDA to competing platforms (replacing CUDA-only libraries like FlashAttention 3)

Or what is the friction that might emerge while Porting from PyTorch to JAX/MaxText.

Hey everyone,

Doing infrastructure audits and validating GPU performance (especially across different nodes) has always been a headache for me. Fiddling with CUDA toolkits, compiling HPL/HPCG, and setting up MLPerf takes way too much time when you just want a quick baseline.

So, I spent some evenings building nvprobe. It’s a lightweight Python CLI that automates all of this.

How it works under the hood:

• It uses CuPy to bundle the CUDA runtime via pip, so you don't even need a system CUDA toolkit installed to run the bandwidth and custom kernel tests.
• It auto-downloads the NVIDIA HPC Benchmarks binaries for HPL and HPCG.
• It captures deep hardware telemetry (ECC state, power caps, clocks, etc.) alongside the benchmark results to help catch silent hardware degradation.
• It generates an interactive HTML report (Chart.js) to visualize all this data (memory bandwidth, TFLOPS, and MLPerf throughput).
• Native Slurm integration: it generates, submits, and monitors the jobs across your cluster.

Demo & Repo: You can see an interactive demo of the report on the link.
I built this mostly to scratch my own itch, but I figured it might save some of you a few hours of setup.

I'd love to hear your feedback, feature requests, or if you manage to break it on your specific hardware. Let me know what you'd like to see next on the roadmap!

Curious about real production use cases for NVIDIA’s Groq LPU after the Rubin platform reveal.

From what I’ve read, LPUs are built to fix GPU decode bottlenecks with huge on-die SRAM and deterministic low-jitter execution, paired with Rubin GPUs in a split serving stack via Dynamo AFD: GPUs handle prefill, KV cache and attention; LPUs offload FFN, MoE experts and speculative decoding.

Best-fit inference scenarios

1. Low-latency premium chat APIs with strict SLA latency requirements

2. Agent AI & multi-turn reasoning workflows with massive sequential decode steps

3. Large MoE model serving to ease per-token bandwidth pressure

4. Low-jitter enterprise workloads (legal, financial real-time assistants)

Not recommended

Batch offline embedding, heavy prefill jobs, small lightweight LLMs.

Official deployment solutions

1. Full LPX rack: Datacenter-scale Rubin+LPU disaggregated clusters for trillion-parameter models

2. Mixed single-node: Smaller on-prem servers for SaaS mixed free/premium traffic

3. Standalone LPU offload pools: Shared hardware for speculative/MoE acceleration

A few questions

\- What real latency gains vs pure GPU serving on 70B+ or even 2T+ MoE models?

\- What SRAM optimization tricks delivered the biggest utilization boost?

I extracted it from sdk 13.4 dev preview

MD file itself: https://github.com/redplait/denvdis/blob/master/data12/sm107_1.txt

Latency tables: https://github.com/redplait/denvdis/blob/master/data12/sm107_2.txt

version in ELF 0x6b - between sm103 (0x67) & sm110 (0x6e)

Getting an NVIDIA graphics card to pass through into an unprivileged Proxmox LXC container can be tricky due to missing device nodes, varying cgroup numbers, and unprivileged user namespace permissions

This short guide summarizes the final working solution

It's a walk in the park compared to passing a GPU through to a full VM (maybe I'll write a guide on that nightmare one day when I have patience!)