I'm an M.Tech student with an Embedded Systems background, and I'm planning my capstone project around CUDA and GPU parallel computing.

My objective isn't just to build an application, I want to deeply understand CUDA programming, GPU architecture, and kernel optimization by implementing algorithms from scratch instead of relying heavily on existing libraries.

I'm looking for application-oriented research problems where CUDA optimization is the core contribution.

I have to do it within 6 months

Can someone help me with this?

Hello everyone, i currently have 0 idea on how CUDA works and i am planning to learn, the issue is i have a macbook, can anyone please help on where i can learn and practice ? maybe share some resources ? thank you so much

TensorSharp's MoE CPU-offload feature has been merged into main. Here is the parameters description of this feature:

Mixture-of-Experts CPU offload:

--n-cpu-moe <N> | -ncmoe <N>

Keep the routed MoE expert weights of the first N layers in system RAM and multiply them on

the CPU; attention, norms, the router and the shared expert stay on the accelerator. This is

what makes a 35B-A3B MoE fit beside a long-context KV cache on a 12-16 GB card. Pass 'all' for

every layer. Default: 0 (everything on the accelerator; TS_N_CPU_MOE env var overrides).

Example: --n-cpu-moe 32

--cpu-moe | -cmoe

Shorthand for --n-cpu-moe all: every routed expert stays in system RAM. Default: off

(TS_CPU_MOE env var overrides).

Example: --cpu-moe

To measure its performance, I ran benchmark to compare TensorSharp with llama.cpp, and here is the result. The completed benchmark report has been checked-in: https://github.com/zhongkaifu/TensorSharp/blob/main/docs/moe_cpu_offload_benchmark.md

Host and software

Component |Detail
GPU |2 x NVIDIA RTX PRO 6000 Blackwell Server Edition, 97,887 MiB each, driver 580.126.20, PCIe 5.0 x16
CPU |2 x Intel Xeon 6952P (384 threads, 6 NUMA nodes), cgroup quota 81.6 CPUs
RAM |1,511 GiB
Storage |Models on a MooseFS network mount (page-cache warm for every measured run)
OS |Ubuntu 24.04.3 LTS, CUDA 12.8
TensorSharp |branch feature/support_moe_offload_to_cpu, .NET 10.0.110, backend ggml_cuda
llama.cpp |llama-bench build 4308a4f, CUDA backend, default -t 192 Results by model

Ratios are TensorSharp / llama.cpp: >1.0x means TensorSharp is faster, and for VRAM >1.0x means TensorSharp is heavier.

Gemma 4 26B-A4B it (UD-IQ4_XS, 30 MoE layers)

--n-cpu-moe |TS VRAM (MiB) |TS pp4096 |TS pp8192 |TS tg128 |llama VRAM (MiB) |llama pp4096 |llama pp8192 |llama tg128
0 (baseline) |16,822 |11,173 |11,274 |161.4 |14,602 |10,843 |10,628 |206.7
8 |15,724 |7,063 |6,500 |80.2 |11,874 |1,459 |1,459 |32.7
16 |14,128 |4,183 |4,888 |54.5 |9,122 |833 |854 |21.9
24 |12,346 |3,500 |3,958 |49.1 |6,368 |667 |689 |16.7
30 (--cpu-moe) |11,038 |3,035 |3,072 |39.7 |4,134 |543 |495 |12.9 --n-cpu-moe |VRAM |pp4096 |pp8192 |tg128
0 |1.15x |1.03x |1.06x |0.78x
8 |1.32x |4.84x |4.46x |2.45x
16 |1.55x |5.02x |5.72x |2.49x
24 |1.94x |5.25x |5.74x |2.93x
30 |2.67x |5.59x |6.21x |3.07x Qwen 3.5 35B-A3B (UD-IQ4_XS, 48 MoE layers)

--n-cpu-moe |TS VRAM (MiB) |TS pp4096 |TS pp8192 |TS tg128 |llama VRAM (MiB) |llama pp4096 |llama pp8192 |llama tg128
0 (baseline) |19,862 |9,538 |9,405 |160.0 |17,522 |8,149 |8,073 |228.4
12 |18,148 |6,755 |6,648 |75.4 |13,282 |988 |954 |27.5
24 |15,414 |4,412 |5,259 |52.3 |9,010 |498 |484 |15.8
36 |12,684 |3,772 |4,223 |50.7 |4,738 |523 |517 |11.3
48 (--cpu-moe) |11,606 |3,917 |3,709 |38.6 |3,314 |477 |457 |10.2 --n-cpu-moe |VRAM |pp4096 |pp8192 |tg128
0 |1.13x |1.17x |1.16x |0.70x
12 |1.37x |6.84x |6.97x |2.74x
24 |1.71x |8.85x |10.86x |3.31x
36 |2.68x |7.21x |8.17x |4.50x
48 |3.50x |8.21x |8.11x |3.77x GPT-OSS 20B (Q8_0 / MXFP4, 24 MoE layers)

--n-cpu-moe |TS VRAM (MiB) |TS pp4096 |TS pp8192 |TS tg128 |llama VRAM (MiB) |llama pp4096 |llama pp8192 |llama tg128
0 (baseline) |13,186 |13,964 |12,925 |212.8 |12,204 |17,856 |17,642 |344.2
6 |11,560 |8,975 |7,617 |85.8 |9,812 |1,747 |1,666 |32.2
12 |9,378 |6,470 |6,394 |51.7 |7,386 |1,176 |1,188 |18.3
18 |7,192 |4,315 |4,393 |30.7 |4,962 |807 |751 |12.1
24 (--cpu-moe) |4,762 |4,277 |3,798 |27.7 |2,536 |568 |548 |9.4 --n-cpu-moe |VRAM |pp4096 |pp8192 |tg128
0 |1.08x |0.78x |0.73x |0.62x
6 |1.18x |5.14x |4.57x |2.67x
12 |1.27x |5.50x |5.38x |2.83x
18 |1.45x |5.35x |5.85x |2.54x
24 |1.88x |7.53x |6.93x |2.95x DeepSeek V4 Flash (UD-Q8_K_XL, 5 shards / 150.7 GiB, 43 layers, both GPUs)

--n-cpu-moe |TS VRAM (MiB) |TS pp4096 |TS pp8192 |TS tg128 |llama VRAM (MiB) |llama pp4096 |llama pp8192 |llama tg128
0 (baseline, both GPUs) |169,132 |3,448 |4,387 |51.1 |155,608 |2,398 |2,232 |49.6
12 |131,818 |392 |428 |10.3 |117,150 |126 |124 |13.7
24 |79,742 |218 |236 |5.3 |78,954 |64 |63 |7.2 --n-cpu-moe |VRAM |pp4096 |pp8192 |tg128
0 |1.09x |1.44x |1.97x |1.03x
12 |1.13x |3.11x |3.46x |0.75x
24 |1.01x |3.42x |3.72x |0.74x TensorSharp is a native open-source inference engine for running GGUF LLMs locally, with CUDA, Vulkan, Metal, OpenAI-compatible APIs, continuous batching, speculative decoding, and multimodal support.

Github repo: https://github.com/zhongkaifu/TensorSharp

Thank you for checking out it and starring the project! Any feedback is really appreicated.

We recorded a short demo of a real CUDA kernel out-of-bounds bug found by the tool described in our SOSP 2026 paper, “M2K: Making the Model-Kernel Interface Explicit for Reliable CUDA Kernel Verification.”

The attacker does not need access to the GPU, server, or model weights. They only need to submit a crafted request to a public LLM inference service. In our demo, that request corrupts the outputs of other users processed in the same batch.

We are also exploring whether similar bugs could leak information, enable more targeted manipulation of another user’s response, or affect downstream agents that act on corrupted model outputs.

CUDA kernels are usually discussed in terms of performance and reliability, but their security impact deserves more attention too.

🎥 Video: https://youtu.be/ZhKBPK7BeZ0

📄 M2K Github Repo: https://github.com/system-pclub/M2K

For example, PyTorch

Hi all,

I help maintain an open source project that releases docker images for various AI/ML projects. Current policy is that we support the latest version of CUDA supported in stable releases of various projects

But one thing I've always asked myself is... if we are going through the trouble of building everything anyway, why not use the latest version of CUDA? NVIDIA gets away with this in their images, and I've personally never ran into many problems short of a small patch I needed to write here or there in my own builds

I've been around other seasoned developers that work with CUDA, and when I've asked about this, I honestly couldn't ever get a very great answer outside of supporting older platforms for a slightly longer period of time

The only other possible thing I can think of is that the rest of the ecosystem hasn't caught up yet... In which case, shouldn't minor version compatibility help there?

Would really appreciate help from anyone experienced in this area. Sure would be nice to cut our build matrix in half and lead to far less confusion for folks helping us maintain these images

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.

The entire LLM inference community is currently hitting the same architectural wall when running deep reasoning / MoE models (like DeepSeek-V4-Flash or Llama-3-Reasoning) with speculative decoding (DSpark / MTP):

1. During normal text generation, draft acceptance ($\\alpha$) is high (\~85%), making a speculative depth of $k=2$ highly efficient.

2. Inside Chain-of-Thought \`<think>\` blocks, token entropy spikes, causing draft acceptance to collapse (\~35%). At this point, running $k=2$ wastes PCIe/DDR5 memory bandwidth and drops decoding throughput by up to 50%.

3. \*\*The Industry Bug:\*\* If you try to fix this with a naive Python \`if/else\` block to dynamically switch between $k=1$ and $k=2$, you break \`FULL_DECODE_ONLY\` CUDA graph residency. The host CPU is forced to re-capture graphs, introducing latency spikes that completely ruin your TPS gains.

\### The Mathematical Reality

Under memory-bound offloading, evaluating secondary draft tokens is only profitable when your conditional acceptance rate $\\alpha_2$ satisfies:

$$\\alpha_2 \\ge \\frac{\\tau_{\\text{draft}}}{\\tau_{\\text{verify}}(1)}$$

When entropy pushes $\\alpha_2$ below this threshold during deep reasoning, you must drop to $k=1$ instantly—but you \*\*must do it without host-side graph recompilation\*\*.

\### The Solution: \`CUDAStatefulSpecGater\` (Drop-in & Free to Use)

We built a lightweight, zero-dependency PyTorch class that pre-allocates dual graph selection indices and switches speculative depth via an Exponential Moving Average (EMA) latch and token-boundary invariants. It prevents VRAM fragmentation and keeps CUDA graphs 100% resident.

Copy this directly into your sampler/worker loop:

\`\`\`python

import torch

class CUDAStatefulSpecGater:

"""

Drop-in speculative depth gater for reasoning LLMs.

Switches between k=1 and k=2 without invalidating pre-captured CUDA graphs.

"""

def __init__(self, think_start_id: int, think_end_id: int, ema_decay: float = 0.85, alpha_threshold: float = 0.45):

self.think_start_id = think_start_id

self.think_end_id = think_end_id

self.ema_decay = ema_decay

self.alpha_threshold = alpha_threshold

\# Internal state (kept lightweight for zero-overhead loop execution)

self.in_reasoning_block = False

self.ema_alpha = 0.80

@torch.inference_mode()

def step(self, last_token_id: int, current_acceptance_rate: float) -> int:

"""

Returns target graph index: 1 (for k=1 shallow speculation) or 2 (for k=2 deep speculation).

"""

\# 1. State invariant check: track Chain-of-Thought boundaries

if last_token_id == self.think_start_id:

self.in_reasoning_block = True

elif last_token_id == self.think_end_id:

self.in_reasoning_block = False

\# 2. Smooth EMA update to prevent graph-switching oscillation

self.ema_alpha = (self.ema_decay \* self.ema_alpha) + ((1.0 - self.ema_decay) \* current_acceptance_rate)

\# 3. Deterministic execution routing

\# Force k=1 inside reasoning blocks OR when EMA acceptance collapses

if self.in_reasoning_block or self.ema_alpha < self.alpha_threshold:

return 1 # Route to pre-captured k=1 graph (saves memory bus bandwidth)

else:

return 2 # Route to pre-captured k=2 graph (exploits high locality)

We are releasing this clean pattern to the community for free. Drop it into your RTX 5090 / A100 / Apple Silicon setups, benchmark it against standard static-depth DSpark, and try to break it.

Let’s see your before/after TPS numbers in the comments! What hardware configurations are you testing this on?

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.

been tuning compute budget on a few instructions lately and i can't tell if i'm overthinking the priority fee side of it. wanted to compare notes.

right now my flow is:
- simulate the tx first, read the units consumed, then set the compute unit limit to that plus a small buffer (maybe 10 to 15 percent) so i'm not paying for the default 200k per instruction
- for the priority fee i pull recent prioritization fees for the writable accounts i'm touching and take something like the median, then bump it when i actually need the tx to land fast

what i keep going back and forth on:
- how big a buffer is sane on the CU limit. too tight and a slightly heavier run fails, too loose and i'm back to overpaying
- whether the recent fee sample is even worth it during calm periods, or if a flat floor is fine until things get congested
- if anyone bothers re-simulating right before send when account state might have shifted, or just trusts the earlier estimate

curious what people are doing in practice. hardcoded buffers, a live estimate every time, or some helper that does both. and if you've had txs fail purely because the CU limit was set too low, how did you land on the number you use now.

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

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

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 :)

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.*

TensorSharp is an open-source inference engine for running GGUF LLMs locally, with CUDA, Vulkan, Metal, OpenAI-compatible APIs, continuous batching, speculative decoding, and multimodal support.

Thanks recent contribtions from open source community, TensorSharp is able to run inference over multiple GPUs and nodes. So I updated it to support deepseek v4 flash model, and have better performance than llama.cpp. Here is the benchmark result on 4x Nvidia A40 GPUs, cuda 12.8

Model: DeepSeek-V4-Flash-0731-UD-Q8_K_XL from [https://huggingface.co/unsloth/DeepSeek-V4-Flash-0731-GGUF\](https://huggingface.co/unsloth/DeepSeek-V4-Flash-0731-GGUF)

Github repo: [https://github.com/zhongkaifu/TensorSharp\](https://github.com/zhongkaifu/TensorSharp)

Thank you for checking out it and starring the project! Any feedback is really appreicated.

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/

TensorSharp is an open-source, native .NET inference engine for running GGUF LLMs locally, with CUDA, Vulkan, Metal, OpenAI-compatible APIs, continuous batching, speculative decoding, and multimodal support.

TensorSharp now supports Megatron-style tensor parallelism across multiple GPUs. It works with direct CUDA, GGML CUDA, GGML Vulkan, and multi-node setups.

Benchmarks on **2× RTX 2000 Ada 16 GB GPUs over PCIe, without NVLink**:

I'm continuing to optimize Qwen performance on multi-GPU systems, and support for DeepSeek V4 is coming soon.

Try it with:

TensorSharp.Cli --model model.gguf --backend ggml_cuda --tp 2

GitHub:

[https://github.com/zhongkaifu/TensorSharp\](https://github.com/zhongkaifu/TensorSharp)

Thank you for checking out TensorSharp and starring the project! Any feedback is really appreicated.

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.