I'd prefer a layout with a larger block size that doesn't require repacking.

Hi @jeffbolznv

Thanks for the suggestion, I went back at it and now have a slightly faster implementation, but without using any repack. It's still keeping my preloaded tile idea, which is the fastest, and we can put those right into the gguf and not really an issue. 50+ variations of CUDA trials running all day found best tile was just doubling the layout of the previous tile. (so now 832B). I've got CPU rewritten to use that directly now too. On CUDA, with the bigger block and some new tweaks, it gained about 5% for both pp512/tg128 Qwen3.5-4b. the model loads ~1.5s faster without needing the repack. About 18,200 pp512, 191tg. Qwen3.6-35A-A3B is about 7500/170.
I'll post the updated commit and upload the new GGUF on HF after clean up and build checks, and verifying ppl/kld runs all come back correct.

I'd prefer a layout with a larger block size that doesn't require repacking.

I don't thing you can have high perf without repacking. And the block size is really related to hardware, tensor sizes.
So for me keep gguf as simple as possible, and let the backend do is optimised repacking.

@Djip007 interestingly enough, the newer "tile-in-gguf" layout on cpu is also faster:
16.5/pp512 and 6.0/tg128. The original was about half that. Was not really focusing primarily about CPU speed however.

16.5/pp512 and 6.0/tg128

Didn't know what CPU you have, and if it is for the Qwen3.5-4b. but for me look really slow.

_mm512_dpwssd_epi32 can use dot2 on signed int16: so base bloc is A[M/M0][K/K0][M0=16][K0=2]
other use dot4 on int8 (mix signed/unsigned so more complicated...) bloc is A[M/M0][K/K0][M0=16][K0=4]

this are for AVX512...

Pushed a new refactored version. No more block layout and no more repack into tiles.
Now it is storing the fast tile layout directly in the gguf, and the CPU kernel consumes the tile version directly (it's faster, too). Re-quantized the model again and have better results, too:

Mean PPL(Q)                   :   9.998305 ±   0.071790
Mean PPL(base)                :   9.943929 ±   0.071459
Cor(ln(PPL(Q)), ln(PPL(base))):  99.73%
Mean ln(PPL(Q)/PPL(base))     :   0.005453 ±   0.000528
Mean PPL(Q)/PPL(base)         :   1.005468 ±   0.000531
Mean PPL(Q)-PPL(base)         :   0.054377 ±   0.005275

====== KL divergence statistics ======
Mean    KLD:   0.015374 ±   0.000300
Maximum KLD:  19.809378
Median  KLD:   0.008529

RMS Δp    :  3.437 ± 0.044 %
Same top p: 93.487 ± 0.064 %

The weight and input scale are included in the tile header as well, so there's no more code plumbing needed to load up separate tensors or deal with any derived tensor pathway. They are just easily factored in directly in the vecdot.

nice!

did you manage to store the

    float         weight_scale;
    float         input_scale;
    const float * weight_scales;
    const float * input_scales;

in the gguf file? or have an idea for that?
at the begining/end of the tensor data?

(I think of what we can do for FP8 ... )

nice!

did you manage to store the

    float         weight_scale;
    float         input_scale;
    const float * weight_scales;
    const float * input_scales;

in the gguf file? or have an idea for that? at the begining/end of the tensor data?

(I think of what we can do for FP8 ... )

Yes
This is stored in the gguf file and is part of the outer tile. It is faster this way. Thephysical GGUF physical layout is a derived CUDA tiled serialization of the MXFP6 data.
The quantizer writes those header fields directly into the tensor payload, so I stopped it from emitting separate MXFP6 .scale or .input_scale tensors which just cause more code work.
Yes, this goes for MXFP8 too.. I already have MXFP8 and FP8 working; I've been doing the code in parallel, I just don't want to flood too much here so we do one thing at a time. MXFP6 activations are using FP8 with just a very small amount of code to bypass Q8 and get to f4f6f8 which is faster than MXFP6 x Q8. But let's do one thing at a time :)

Yes, this goes for MXFP8 too

I thing at native E4M3 models, but can be good too. (I had some CPU/RDNA3 kernel for that, but did not finish (start) the gguf file. I "only" had some on load convertion from BF16 model to packed E4M3 / E3M4 , with different scaling schema.
I only did not find (understand) how the gguf data is construct/read in this PR.

But let's do one thing at a time :)

Yes! 👍

I only did not find (understand) how the gguf data is construct/read in this PR.

Hope this (generated image) helps..

For the weight_scales / input_scales
You think to store it at the end of the tiles. ?

No, they actually are at the front/top, the header tensor_mxfp6 contains tiles[],

struct tensor_mxfp6 {
    float         weight_scale;
    float         input_scale;
    const float * weight_scales; <--- NULL pointer in file
    const float * input_scales; <-- NULL pointer  in file
    tile_mxfp6   tiles[];
};

tile_mxfp6 is the container type for these tiles[], underneath tensor_mxfp6:

tensor_mxfp6
   (tensor weight/input scales here)
   tile_mxfp6 tiles[]
      tile_mxfp6[0]
         frag[0]: tile_mxfp6_frag is the container type
         frag[1]
      tile_mxfp6[1]
      tile_mxfp6[2]   

One tile_mxfp6_frag:

   tile_mxfp6_frag:
      lane[32][3]  32 lanes,  each has 3 mxfp6 payloads (weights)
      scale[32]   32 MX scales (the subblock scales, E8M0)

To clarify just a bit more about why this is here:

    const float * weight_scales; <--- NULL pointer in file
    const float * input_scales; <-- NULL pointer  in file

These are placeholders in the GGUF, but not actively used in the CPU version, so it appears it doesn't do anything now. They get filled on the CUDA version only. MUL_MAT_ID will select the channel_x expert scale from inside the kernel. That way it does weight_scales[channel_x] in the vecdot instead of launching separate get_rows + mul. That increases both speed and keeps the accuracy applying the weight scale without losing any precision (eg, as derived). CPU side is fine without patching because it keeps the existing separate scale nodes in the graph. On CUDA it will nullify the graph .scale node so they don't get late scaled and get applied directly:

1. first attach expert scale tensor into weight->src[0]
2. set to null the graph .scale node
3. matmul_id gets correct expert channel id
4. vecdot does weight_scales[channel_x] inside the kernel

keeps the existing separate scale nodes in the graph

Ho. That what I do not like (and gg used not too) because this mean model graph depend on quantized type used (ie: all model have to be update).
So I'll look if we can add it at the end of the data(?). (with for exemple weight_scales / input_scales as delta position from data/tiles.

but it is more for #22042 or some other Discussions

1 point : be careful with this header on mmap file... I don't know if you can write it. (for CPU tensors or similare)

keeps the existing separate scale nodes in the graph

Ho. That what I do not like (and gg used not too) because this mean model graph depend on quantized type used (ie: all model have to be update). So I'll look if we can add it at the end of the data(?). (with for exemple weight_scales / input_scales as delta position from data/tiles.

but it is more for #22042 or some other Discussions

1 point : be careful with this header on mmap file... I don't know if you can write it. (for CPU tensors or similare)

We can certainly make the CPU side do the same as CUDA, it already handles the mmap issue and does much of the math (maybe as you think we could improve this further, as far as coming up with the best GGUF layout). Perhaps I need clean up my CUDA MXPF6 side and post that code to a fork for relevance, but this how I handle that there:

static bool ggml_cuda_set_tensor_mxfp6(ggml_tensor * tensor, const void * data, size_t offset, size_t size, int device) {

    ggml_tensor * storage = tensor->view_src ? tensor->view_src : tensor;
    const size_t logical_size = ggml_nbytes(storage);
    const size_t packed_size = ggml_cuda_mxfp6_e2m3_tensor_alloc_size(storage);
    GGML_ASSERT(logical_size == packed_size);
    const size_t storage_offset = tensor->view_src ? tensor->view_offs : 0;

    CUDA_CHECK(cudaMemcpyAsync((char *) storage->data + storage_offset + offset, data, size, cudaMemcpyHostToDevice, cudaStreamPerThread));
    if (storage_offset == 0 && offset == 0 && size >= MXFP6_HEADER_OFFSET) {
        tile_mxfp6 header;
        memcpy(&header, data, MXFP6_HEADER_OFFSET);
        ggml_cuda_mxfp6_e2m3_patch_tensor_header(storage, &header);
        CUDA_CHECK(cudaMemcpyAsync(storage->data, &header, MXFP6_HEADER_OFFSET, cudaMemcpyHostToDevice, cudaStreamPerThread));
    }
    CUDA_CHECK(cudaStreamSynchronize(cudaStreamPerThread));
    return true;
}
static inline void ggml_cuda_mxfp6_e2m3_patch_tensor_header(const ggml_tensor * tensor, tile_mxfp6_e2m3_blackwell_tensor * dst) {
    const ggml_tensor * weight_scale_t = tensor->src[0];
    const ggml_tensor * input_scale_t  = tensor->src[1];

    dst->weight_scales = weight_scale_t != nullptr && !ggml_is_scalar(weight_scale_t) &&
        weight_scale_t->data != nullptr && weight_scale_t->buffer != nullptr &&
        !ggml_backend_buffer_is_host(weight_scale_t->buffer) ? (const float *) weight_scale_t->data : nullptr;
    dst->input_scales = input_scale_t != nullptr && !ggml_is_scalar(input_scale_t) &&
        input_scale_t->data != nullptr && input_scale_t->buffer != nullptr &&
        !ggml_backend_buffer_is_host(input_scale_t->buffer) ? (const float *) input_scale_t->data : nullptr;
}
auto get_mapped_scalar_f32 = [&ml](ggml_tensor * tensor, float & value) {
    if (tensor == nullptr || !ggml_is_scalar(tensor) || tensor->type != GGML_TYPE_F32) {
        return false;
    }
    const auto * weight = ml.get_weight(ggml_get_name(tensor));
    if (weight == nullptr || weight->idx >= ml.mappings.size()) {
        return false;
    }
    const auto & mapping = ml.mappings.at(weight->idx);
    memcpy(&value, (const uint8_t *) mapping->addr() + weight->offs, sizeof(value));
    return true;
};

}

    const ggml_tensor * weight_scale_t = tensor->src[0];
    const ggml_tensor * input_scale_t  = tensor->src[1];

Is the scale define as source of the weight tensor?
That may be for the ggml-runtime/backend the more interesting !
This will solve many question (like the type, sizes ...)

    const ggml_tensor * weight_scale_t = tensor->src[0];
    const ggml_tensor * input_scale_t  = tensor->src[1];

Is the scale define as source of the weight tensor? That may be for the ggml-runtime/backend the more interesting ! This will solve many question (like the type, sizes ...)

Yes!

weight->src[0] = weight_scale;
weight->src[1] = input_scale;

and

const ggml_tensor * src0 = dst->src[0];        // weight tensor
const ggml_tensor * scale = src0->src[0];      // weight scale

This does not mess up any of the other models so there is no need to change anything for anything else (unless improvements for future types)

    auto attach_native_blackwell_scales = [](ggml_tensor * weight, ggml_tensor *& weight_scale, ggml_tensor * input_scale, bool use_cuda_native_scales) {
        if (weight == nullptr || (weight->type != GGML_TYPE_NVFP4 && weight->type != GGML_TYPE_MXFP6_E2M3)) {
            return;
        }

        // Files produced by llama-quantize store plain native blocks without an
        // auxiliary per-tensor scale. Treat that as an identity.
        if (weight_scale == nullptr) {
            const float identity_scale = 1.0f;
            memcpy(&weight->op_params[0], &identity_scale, sizeof(identity_scale));
        }

        weight->src[0] = weight_scale; // matmul reads attached scales before CUDA hides graph *_s
        weight->src[1] = input_scale;
        if (use_cuda_native_scales && weight_scale != nullptr) {
            weight_scale = nullptr; // CUDA applies attached scales inside matmul, so graph skips late ggml_mul
        }
    };

**correction, this was when using repack, after the new layout, newer cuda version doesnt need op_params for a placeholder anymore

Yes!

weight->src[0] = weight_scale;
weight->src[1] = input_scale;

and

const ggml_tensor * src0 = dst->src[0];        // weight tensor
const ggml_tensor * scale = src0->src[0];      // weight scale

That's what I missed!
(https://github.com/ggml-org/llama.cpp/blob/master/src/llama-model.cpp#L1279)

Uploaded new MXFP6 models onto HF:
Qwen3.6-35B-A3B-MXFP6
Qwen3.6-27B-MXFP6

To demonstrate how the no-repack tile layout stored in the GGUF can work without big GGML changes, and keeping the scales in the same tile as the CPU version, I posted the CUDA implementation here (not for PR and not clean enough, but as POC):
mxfp6-cuda

This converts q8_1 tiles to FP8 for MMVQ, so doing native Blackwell MXFP6 x FP8 MMA. This makes MXFP6 almost 20% faster than NVFP4 for tg on this model while keeping the extra quality.

  Device 0: NVIDIA GeForce RTX 5090, compute capability 12.0, VMM: yes, VRAM: 32606 MiB
| model                          |       size |     params | backend    | ngl |            test |                  t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | --------------: | -------------------: |
| qwen35moe 35B.A3B MXFP6 - E2M3 |  26.46 GiB |    34.66 B | CUDA       |  99 |           pp512 |      8094.43 ± 49.53 |
| qwen35moe 35B.A3B MXFP6 - E2M3 |  26.46 GiB |    34.66 B | CUDA       |  99 |           tg128 |        188.10 ± 3.20 |

  Device 0: NVIDIA GeForce RTX 5090, compute capability 12.0, VMM: yes, VRAM: 32606 MiB
| model                          |       size |     params | backend    | ngl |            test |                  t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | --------------: | -------------------: |
| qwen35moe 35B.A3B NVFP4        |  21.48 GiB |    34.66 B | CUDA       |  99 |           pp512 |      8220.18 ± 57.89 |
| qwen35moe 35B.A3B NVFP4        |  21.48 GiB |    34.66 B | CUDA       |  99 |           tg128 |        159.53 ± 0.82 |

Hoping for any feedback on making this even better!

The ppl/kld delta on Qwen3.6-27b vs NVFP4 is really striking! MXFP4 vs BF16:

Mean PPL(Q)                   :   6.918624 ±   0.045510
Mean PPL(base)                :   6.900856 ±   0.045374
Cor(ln(PPL(Q)), ln(PPL(base))):  99.52%
Mean ln(PPL(Q)/PPL(base))     :   0.002571 ±   0.000645
Mean PPL(Q)/PPL(base)         :   1.002575 ±   0.000646
Mean PPL(Q)-PPL(base)         :   0.017768 ±   0.00445
Mean    KLD:   0.018925 ±   0.000647
RMS Δp    :  3.662 ± 0.060 %
Same top p: 95.023 ± 0.057 %

and for NVFP4:

====== Perplexity statistics ======
Mean PPL(Q)                   :   7.321749 ±   0.049291
Mean PPL(base)                :   6.900856 ±   0.045374
Cor(ln(PPL(Q)), ln(PPL(base))):  98.00%
Mean ln(PPL(Q)/PPL(base))     :   0.059204 ±   0.001341
Mean PPL(Q)/PPL(base)         :   1.060991 ±   0.001423
Mean PPL(Q)-PPL(base)         :   0.420893 ±   0.010245
Mean    KLD:   0.079692 ±   0.001052
RMS Δp    :  7.764 ± 0.063 %
Same top p: 89.020 ± 0.081 %

Thanks for the excellent MXFP6 work. I was impressed by this PR and the related discussion, so I tried an experimental ROCm implementation
for RDNA4 with an AMD Radeon AI PRO R9700 32GB GPU, with help from Codex.

This is not PR-ready yet, but I wanted to share some early numbers in case they are useful.

Environment:

• GPU: AMD Radeon AI PRO R9700, gfx1201, 32GB
• Backend: ROCm / HIP
• Model: Qwen3.6 27B GGUF variants
• llama-bench: -ngl 99 -r 5
• KV: f16/f16

Speed results:

quant size pp128 pp512 pp1024 tg128
━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━ ━━━━━━━ ━━━━━━━━━━━━ ━━━━━━━━━━━━ ━━━━━━━━━━━━ ━━━━━━━━━━━━━━
MXFP6_E2M3, compact/original layout ~21GB ~618 tok/s ~830 tok/s ~566 tok/s ~24.67 tok/s
─────────────────────────────────────── ─────── ──────────── ──────────── ──────────── ──────────────
MXFP6_E2M3, RDNA4 repacked F26 layout ~21GB ~680 tok/s ~910 tok/s ~589 tok/s ~23.72 tok/s
─────────────────────────────────────── ─────── ──────────── ──────────── ──────────── ──────────────
Q6_K ~21GB ~501 tok/s ~664 tok/s ~484 tok/s ~22.81 tok/s
─────────────────────────────────────── ─────── ──────────── ──────────── ──────────── ──────────────
UD-Q5_K_XL ~19GB ~492 tok/s ~773 tok/s ~509 tok/s ~24.15 tok/s
─────────────────────────────────────── ─────── ──────────── ──────────── ──────────── ──────────────
Q4_K_M ~16GB ~693 tok/s ~959 tok/s ~621 tok/s ~26.39 tok/s

Short WikiText-2 perplexity check, ctx 512, chunks 64:

quant PPL
━━━━━━━━━━━━ ━━━━━━━━━━━━━━━━━━━━
MXFP6_E2M3 6.5708 +/- 0.12733
──────────── ────────────────────
UD-Q5_K_XL 6.5689 +/- 0.12722

A few observations from the RDNA4 side:

• MXFP6 is already faster than Q6_K for both PP and TG on this setup.
• PP/GEMM benefits from an RDNA4-friendly repacked layout.
• TG/matvec did not improve with the RDNA4 repack; the original compact layout was still faster for TG.
• Against UD-Q5_K_XL, MXFP6 has similar short PPL and slightly better TG/PP, but the advantage is not large yet.
• Q4_K_M remains very strong for TG and size.

My current impression is that a hardware-neutral GGUF MXFP6_E2M3 layout plus backend-specific load-time repacking is probably the right
direction. For RDNA4, the main open problem seems to be TG/matvec optimization rather than PP/GEMM.

I will keep experimenting before publishing code, but I wanted to share these early ROCm/RDNA4 data points.

Thanks for the excellent MXFP6 work. I was impressed by this PR and the related discussion, so I tried an experimental ROCm implementation for RDNA4 with an AMD Radeon AI PRO R9700 32GB GPU, with help from Codex.

This is not PR-ready yet, but I wanted to share some early numbers in case they are useful.

Environment:

• GPU: AMD Radeon AI PRO R9700, gfx1201, 32GB
• Backend: ROCm / HIP
• Model: Qwen3.6 27B GGUF variants
• llama-bench: -ngl 99 -r 5
• KV: f16/f16

Speed results:

quant size pp128 pp512 pp1024 tg128 ━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━ ━━━━━━━ ━━━━━━━━━━━━ ━━━━━━━━━━━━ ━━━━━━━━━━━━ ━━━━━━━━━━━━━━ MXFP6_E2M3, compact/original layout ~21GB ~618 tok/s ~830 tok/s ~566 tok/s ~24.67 tok/s ─────────────────────────────────────── ─────── ──────────── ──────────── ──────────── ────────────── MXFP6_E2M3, RDNA4 repacked F26 layout ~21GB ~680 tok/s ~910 tok/s ~589 tok/s ~23.72 tok/s ─────────────────────────────────────── ─────── ──────────── ──────────── ──────────── ────────────── Q6_K ~21GB ~501 tok/s ~664 tok/s ~484 tok/s ~22.81 tok/s ─────────────────────────────────────── ─────── ──────────── ──────────── ──────────── ────────────── UD-Q5_K_XL ~19GB ~492 tok/s ~773 tok/s ~509 tok/s ~24.15 tok/s ─────────────────────────────────────── ─────── ──────────── ──────────── ──────────── ────────────── Q4_K_M ~16GB ~693 tok/s ~959 tok/s ~621 tok/s ~26.39 tok/s

Short WikiText-2 perplexity check, ctx 512, chunks 64:

quant PPL ━━━━━━━━━━━━ ━━━━━━━━━━━━━━━━━━━━ MXFP6_E2M3 6.5708 +/- 0.12733 ──────────── ──────────────────── UD-Q5_K_XL 6.5689 +/- 0.12722

A few observations from the RDNA4 side:

• MXFP6 is already faster than Q6_K for both PP and TG on this setup.
• PP/GEMM benefits from an RDNA4-friendly repacked layout.
• TG/matvec did not improve with the RDNA4 repack; the original compact layout was still faster for TG.
• Against UD-Q5_K_XL, MXFP6 has similar short PPL and slightly better TG/PP, but the advantage is not large yet.
• Q4_K_M remains very strong for TG and size.

My current impression is that a hardware-neutral GGUF MXFP6_E2M3 layout plus backend-specific load-time repacking is probably the right direction. For RDNA4, the main open problem seems to be TG/matvec optimization rather than PP/GEMM.

I will keep experimenting before publishing code, but I wanted to share these early ROCm/RDNA4 data points.

Thanks @kominsoo appreciate the kind words. A few questions:

• Did you include the weight and input tensor scales? They aren't part of the MXFP6 standard but I still found them useful without really affecting speed.
• Did you take a look at the cuda version here?
• And, how did you implement TG? My MXFP6 trick for fast tg was to use FP8, but just with a fast conversion from Q8 and not put all the plumbing in. The other trick that got a lot of speedup was keeping everything 6-bit until the last second and inserting "Just In Time" 0s to get 8-bit when sending off to the GPU(required by CUDA). Keeping 0s slowed things down dragging them around otherwise. But I'm not sure how it works on ROCm.

I am actually working on a post now in the discussion forum to talk about the new NVFP4/MXFP6 quantizer with a new technique, you can check that in a bit. It needs the tensor scale to work best, you can give it a try with advanced-gguf-quantizer. I think the best final use case will be combining NVFP4 and MXFP6 to balance out each other as needed.

Thanks, those are exactly the points I need to check next. I did use your mxfp6-cuda branch as reference, but I want to verify whether my
ROCm path is applying the tensor scales correctly before commenting further. I’ll test that and report back once I have cleaner numbers.

Following up now that I've verified actual generation output, not just bench numbers — and I owe a correction: the ~24.67 tok/s TG for the compact/original
layout in my earlier table was invalid due to an MMVQ addressing bug. After fixing it, the corrected and optimized comparative results are below.

Q: Tensor scales?
Yes. I preserve the 32-byte tensor_mxfp6 header (the weight_scale) and apply tensor_scale × per-block E8M0 inside both the vecdot and MMQ kernels. Dropping
this header originally produced garbage outputs — once preserved, the Wikitext-2 PPL successfully restored to normal/clean levels.

Q: Reference to your CUDA branch?
Yes. Your mxfp6-cuda branch served as the primary reference for the tile layout and scale handling — thank you very much for posting it.

Q: How was TG optimized?
This is where RDNA4 differs from your CUDA results.
* PP/GEMM uses real FP8 WMMA (wmma_f32_16x16x16_fp8_fp8): We convert MXFP6 to FP8 E4M3 at tile-load time, quantize activations to FP8 E4M3, and execute on
the RDNA4 FP8 matrix core. This resolves the PP regression and beats Q6_K by +21% to +29% across all prompt lengths.
* TG/matvec stays dp4a (int8): We found that routing batch-1 matvec through the FP8 WMMA unit is actually slower than dp4a on RDNA4 (due to tile padding
overhead). Instead, we optimized the dp4a path by hoisting the loop-invariant scales and accumulation chaining.
* As a result, the general MXFP6 GGUF is internally repacked on load time (auto-enabled when an RDNA4 device is present), achieving optimal serving speed on every
axis.

Corrected & Optimized Benchmark Numbers (Tested on AMD Radeon AI PRO R9700, gfx1201, ROCm 7.2, -ngl 99 -r 5)

| Quantization Format | PP128 | PP512 | PP1024 | TG128 | Wikitext-2 PPL (c=2048) |
| :--- | :---: | :---: | :---: | :---: | :---: |
| **MXFP6** *(Auto-repack + TG opt)* | **674.8** | **903.8** | **585.1** | **23.39** | **6.654** *(Lower is Better)* |
| **Q6_K** *(Baseline)* | 525.8 | 699.5 | 482.4 | 23.07 | 6.697 |

Note: TG serving speed practically doubles to **47.82 tok/s** (+104%) when utilising MTP self-speculation (`--spec-type draft-mtp --spec-draft-n-max 4`) on the `llama-server`.

Following up now that I've verified actual generation output, not just bench numbers — and I owe a correction: the ~24.67 tok/s TG for the compact/original layout in my earlier table was invalid due to an MMVQ addressing bug. After fixing it, the corrected and optimized comparative results are below.

Q: Tensor scales? Yes. I preserve the 32-byte tensor_mxfp6 header (the weight_scale) and apply tensor_scale × per-block E8M0 inside both the vecdot and MMQ kernels. Dropping this header originally produced garbage outputs — once preserved, the Wikitext-2 PPL successfully restored to normal/clean levels.

Q: Reference to your CUDA branch? Yes. Your mxfp6-cuda branch served as the primary reference for the tile layout and scale handling — thank you very much for posting it.

Q: How was TG optimized? This is where RDNA4 differs from your CUDA results. * PP/GEMM uses real FP8 WMMA (wmma_f32_16x16x16_fp8_fp8): We convert MXFP6 to FP8 E4M3 at tile-load time, quantize activations to FP8 E4M3, and execute on the RDNA4 FP8 matrix core. This resolves the PP regression and beats Q6_K by +21% to +29% across all prompt lengths. * TG/matvec stays dp4a (int8): We found that routing batch-1 matvec through the FP8 WMMA unit is actually slower than dp4a on RDNA4 (due to tile padding overhead). Instead, we optimized the dp4a path by hoisting the loop-invariant scales and accumulation chaining. * As a result, the general MXFP6 GGUF is internally repacked on load time (auto-enabled when an RDNA4 device is present), achieving optimal serving speed on every axis.

Corrected & Optimized Benchmark Numbers (Tested on AMD Radeon AI PRO R9700, gfx1201, ROCm 7.2, -ngl 99 -r 5)

| Quantization Format | PP128 | PP512 | PP1024 | TG128 | Wikitext-2 PPL (c=2048) |
| :--- | :---: | :---: | :---: | :---: | :---: |
| **MXFP6** *(Auto-repack + TG opt)* | **674.8** | **903.8** | **585.1** | **23.39** | **6.654** *(Lower is Better)* |
| **Q6_K** *(Baseline)* | 525.8 | 699.5 | 482.4 | 23.07 | 6.697 |

Note: TG serving speed practically doubles to **47.82 tok/s** (+104%) when utilising MTP self-speculation (`--spec-type draft-mtp --spec-draft-n-max 4`) on the `llama-server`.

That's great! Thanks for this data. I was minutes ago just now working on improving NVFP4 tg speeds on the repack version by trying FP8 as I had done with MXFP6, while you replied. I also found that dp4a Q8 was much better for single token than the tile path, the same as you were seeing here, but got around it by making a row shadow for MMVQ. Maybe that could help for ROCm too:

CUDA MXFP6 current FP8 TG path:
pp512 17427.2
tg128   230.64

CUDA MXFP6 dp4a TG experiment:
pp512 17344.8
tg128   171.67

NVFP4 forced FP8 TG1: ~120 tok/s
NVFP4 old repack Q8 TG1: ~170 tok/s
NVFP4 row-shadow Q8 TG1: ~193.8 tok/s