We unlocked NVFP4 on the DGX Spark: 20% faster than AWQ!

ever tried this for quantization?

__device__ __forceinline__ uint8_t _sw_float_to_e2m1(float x) {
  uint8_t sign = (uint8_t)((__float_as_uint(x) >> 28) & 8u);
  float ax = fabsf(x);

  // 7 comparisons summed — no branch, only predicates
  uint8_t mag = (ax > 0.25f)
              + (ax >= 0.75f)
              + (ax > 1.25f)
              + (ax >= 1.75f)
              + (ax > 2.5f)
              + (ax >= 3.5f)
              + (ax > 5.0f);

  return sign | mag;
}

your (If-else): ~14–21 Taktzyklen pro Thread (im Schnitt 3–4 Branches evaluiert + Divergenz-Penalty von ~2–3x durch Warp-Serialisierung)

this (branchless): ~7–9 Taktzyklen pro Thread (7× FSETP+IADD, voll pipelined, keine Divergenz)

just to notice:

NVFP4 via CUTLASS already works on both SM120 and SM121 — through FlashInfer CUTLASS JIT:

  • SM120 (Spiegel 2): 157.9 tok/s NVFP4 CUTLASS — OOTB

  • SM121 (DGX Spark): 65.0 tok/s NVFP4 CUTLASS — with FLASHINFER_CUDA_ARCH_LIST="12.0a 12.1a"

We have documented and benchmarked this multiple times. FlashInfer JIT compiles for sm_120a, the CUDA runtime selects the correct binary, and the E2M1 conversion runs in hardware (because the JIT code is compiled for SM120, not SM121).

The FlashInfer way does not compete with marlin + int4 autoround.

The Avarok patch “Software E2M1” therefore solves a problem that only exists when using vLLM’s native NVFP4 path instead of FlashInfer JIT. FlashInfer bypasses this elegantly via dual-arch compilation. The native path is very impressive work!

The real question is: Is vLLM’s native CUTLASS path (with your turbo patches) faster than FlashInfer CUTLASS JIT? And is it faster than Marlin with Int4+Autoround. Tests are running….

PS: i saw futher 121 fixed version test in the new code. Keep in mind that rtx pro 6000 is almost identical to dgx in chipset properties. and it has different version.

Hi @tbraun96,

Great, thanks. How are you checking the accuracy of the models to ensure that the speculative method is lossless and does not affect the generated output or the performance level of the vanilla model?

Sorry, bad choice of words, was a direct semantic borrowing from portuguese (my mother tongue) causing semantic loss, what I meant was reproducible/controlled. The words share some meanings between both languages, but not all.

I found something interesting. Could it be that SM121 HAS the instruction in hardware (cvt.rn.satfinite.e2m1x2.f32)? It’s ptxas (the PTX assembler) that rejects it for the target sm_121a — a toolchain restriction, not a hardware deficit? How would FlashInfer CUTLASS 4.3.5 be able to run on dgx without kernel changes (121 vs 120) if its not the same? I dig deeper into but keeping confused for now.

UPDATE:
CUTLASS does true round-to-nearest-even (banker’s rounding) with guard/round/sticky bits (lines 660–673). Avarok’s version does simple round-to-nearest via thresholds. For most values this is identical, but at the exact midpoints (e.g. 0.75f) there can be a 1 ULP difference. But it seems to fit better?

UPDATE:
It seems, that CUTLASS has “already” a fallback in exmy_base.h, so Avarok’s version is a more performant alternative to the proper IEEE fallback, and there is a further optimization possible.

I made a new branch “turbo”:

There is an english version as well.

its bench.py in the vllm marlin 12x repos.

and its not “lossless”. normally speculative should be lossless in comparison to the vanilla. but my test have shown that BF16 drafters reduce the accuracy when used on nvfp4 with flashinfer cutlass.

If you only want to measure output tokens, tons of output tokens, it’s good and very useful if you prompt: “Read out some texts from Shakespeare.”

Big Thanks!!!

When using GadflyII/Qwen3-Coder-Next-NVFP4, official vllm-thor container worked well for thor. But for spark, I had no stable docker image working until now. Some succeeded on loading but would crash during testing. Tried making one manually but had no success on stable testing using flashinfer until now.

Your image seems working stable and well :)

Thank you everyone for the feedback! I will address each person’s questions when I get to it. For now, I am working on ensuring the GitHub repo properly mirrors what I have locally, is cleaned up, and more organized than before. I would say WAIT before performing a deep investigation. I expect to get a PR merged by the end of the day.

Also, in the same PR, I have managed to reduce the memory usage with almost nonzero change in performance. MTP + Qwen3-Next NVFP4 is currently running with 38GB of RAM, down by about 9 GB compared to the initial release!!

PR: Sparse FP4 2:4 GEMV kernel + Docker build fix + CI pipeline by tbraun96 · Pull Request #3 · Avarok-Cybersecurity/dgx-vllm · GitHub

I’ve got this (GadflyII quantization) running, but slow?!…without MTP, only 28 tg/s at queue depth=1, and with MTP, 0 tokens accepted lowers it to 16 tg/s…any options hints

i must say, i have problems building it straight from the repos. the idee is good, so i will implement the optimized quantization of your e2m() method into cutlass 4.3.5 - lets see if this is the major breakthrough or your concepts of prevent-zero-bytes-transfer or pingpong. maybe a mix of both. i see progress.

Yeah, not surprised. As noted earlier, the repo is a WIP that needs to accurately reflect the code I have locally. I just merged in the PR into main that should make building work (I verified locally — it builds and runs). Can you give it a shot, flash? I want to get to the point where we developers can easily reproduce, tweak, and improve the image iteratively. There’s lots of room for optimization still.

The repo’s “all levels of the stack are ready for dissection” philosophy/strategy will make us fly fast: that means no more waiting for Nvidia, no more waiting for low-level developers tied up elsewhere; we are on our own. Mind you, there should be a PR later that changes sed patches into PR branches as git submodules for each level of the stack. The current approach for building the image is not clean. This is an easy change once we get into a more stable state.

An example of iteration:

I’m currently running micro-optimizations targeting the custom sparse FP4 CUDA kernel. I make sure, before a full build+bench iteration, there will be a substantive change in performance.

the tricky thing is this family/arch/platform build flag that all of the compiler and f.i. flashinfer (not topic here but same problem) check. so i did everything with patches. when there is no upstream into the larger projects, they will fail in near future. i tests your implementations with cutlass 4.3 and it was a mess, cutlass 4.2 where yours based on works but the compiler strikes anyway.

so its more finding the right entrypoint for a long living patch. there is no support in hot fixes.

therefore i suggest using the running framework progress. cutlass has a working fallback for nvfp4. its ieee conform but its slow as hell. it reads like a “sorry for nvfp4 marketing of this dgx thing, but we have to focus on our 100 billion datacenter deals with oracle, openai … ”. so this is an offer from the big money that you’re allowed to plugin. just to use nvfp4 that has a lot of disadvantages to int4 autoround. and… it is nvidia propertiary. keep this in mind.

i’ve read your medium article. its rather a publishing than a blog. a lot of the assumptions in the blog are outdated and hides some very good ideas of you. we should focus.

i only need this dgx works faster, because I bought two of dgx. mission not yet completed.

Gotcha, those patches don’t work because I am using cached/pegged revisions locally that likely are incompatible when you try building it. Yes, that’s a mess. I’m going to make sure anything that requires patches is also tied to a specific revision. Message me, I’ll keep you up to date and we can get the GitHub reproducible for developers.

Yep, we’re basically tier 3 users to Nvidia (tier 2 = RTX 6000, tier 1 = data center GPUs). So, as you point out, we are lower priority.

Oddly, when running some of the nvfp4 Qwen3 Coder Next, I get no speed boost at all. Actually, it is slower fp8 at 23 tok/s. Zero acceptance rate for the spec dec, so I’m guessing I’m just adding pure overhead at that point. Without it, I get 38 tok/s.

I also see this in the logs with coder:

(EngineCore_DP0 pid=94) /app/vllm/vllm/model_executor/layers/fla/ops/utils.py:113: UserWarning: Input tensor shape suggests potential format mismatch: seq_len (9) < num_heads (32). This may indicate the inputs were passed in head-first format [B, H, T, ...] when head_first=False was specified. Please verify your input tensor format matches the expected shape [B, T, H, ...].

Thanks for trying this out on another model. I’d like to give a few pointers.

First, when testing a new model/image, never start right away with speculative decoding. Spec decoding requires:

Proper configuration (requires finding a “sweet spot” of num_speculative_tokens for it to be effective)
Proper Head/Draft models (improper models, or undertrained ones, can degrade performance catastrophically)
Other context-dependent considerations

So, what I’m hearing is that, in the absence of speculative decoding, you’re getting 38 tok/s for the qwen3-next coding variant? I’d be interested in somebody comparing it to the corresponding 4-bit AWQ model.

But certainly, keep speculative decoding off for now. A zero acceptance rate is not useful.

Any Docker images? :D

Yep! The docker image is avarok/dgx-vllm-nvfp4-kernel:v22. See:

Hey everyone! I’ve been away from this for a couple months now - last ran minimax-m2.1-AWQ through Eugr’s repo and honestly I’m a bit scared to dive back in. It pulls you into an infinite rabbit hole 😅 If someone could share the best command to run minimax-m2.5-nvfp4 (Marlin seems to be go-to now, or maybe that CUTLASS patch from flash333), I’ll buy you a week’s worth of coffee as a thank you!

I’m working on it! Testing the nvfp4 REAP version so it can potentially fit everything (+KV cache) into a single DGX. It seems to have auto selected the specialized Marlin NVFP4 MoE Backend for this model, so it should benefit from this image’s speed. Let’s see what happens. I’ll get back to you when complete!