Force NVFP4 W4A8 path for NVFP4_W4A16 layers on Blackwell, where NVFP4 normally uses the native W4A4 path. #24364

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged

ggerganov merged 14 commits into ggml-org:master from ynankani:ynankani/Force_W4A16_NVFP4_to_W4A8

Sep 25, 2026

+501 −153

conversion/base.py

-Original file line number
+Diff line change
@@ Expand Up / @@ -170,6 +170,9 @@ @@
             self.dir_model_card = dir_model  # overridden in convert_lora_to_gguf.py
             self._is_nvfp4 = False
             self._is_mxfp4 = False
+            self._nvfp4_global_algo: str | None = None # checkpoint-wide NVFP4 quant_algo
+            self._nvfp4_layer_algo: dict[str, str | None] = {} # per-layer quant_algo, keyed by HF module path
+            self._prec_a4: dict[str, bool] = {} # gguf tensor name -> can use 4-bit (A4) activations
             self._fp8_as_q8 = fp8_as_q8
             self._fp8_dequantized: set[str] = set()
@@ Expand Down Expand Up / @@ -251,7 +254,7 @@ @@
                         if weight_map is None or not isinstance(weight_map, dict):
                             raise ValueError(f"Can't load 'weight_map' from {index_name!r}")
                         tensor_names_from_index.update(weight_map.keys())
                         part_dict: dict[str, None] = dict.fromkeys(weight_map.values(), None) # ty: ignore[invalid-assignment]
                         part_names = sorted(part_dict.keys())
                 else:
                     weight_map = {}
@@ Expand Down Expand Up / @@ -664,6 +667,18 @@ @@
                     if bias_types:
                         self._fusable_qkv_bias_layers.add(bid)
+        def _tag_prec_a4(self, hf_name: str, gguf_name: str) -> None:
+            # W4A16_NVFP4 should not use 4-bit activations
+            name = hf_name.removesuffix(".weight").removesuffix(".bias")
+            algo = self._nvfp4_global_algo
+            while name:
+                if name in self._nvfp4_layer_algo:
+                    algo = self._nvfp4_layer_algo[name]
+                    break
+                name = name.rpartition(".")[0]
+            if algo == "W4A16_NVFP4":
+                self._prec_a4[gguf_name] = False
         def set_gguf_parameters(self):
             raise NotImplementedError("set_gguf_parameters() must be implemented in subclasses")
@@ Expand Down Expand Up / @@ -837,6 +852,7 @@ @@
             raw, shape = self._nvfp4_pack(weight, scale)
             logger.info(f"Repacked {new_name} with shape {shape} and quantization NVFP4")
             self.gguf_writer.add_tensor(new_name, raw, raw_dtype=gguf.GGMLQuantizationType.NVFP4)
+            self._tag_prec_a4(name, new_name)
             self._write_scale_tensor(new_name.replace(".weight", ".scale"), scale2)
             self._write_scale_tensor(new_name.replace(".weight", ".input_scale"), input_scale)
@@ Expand Down Expand Up / @@ -929,6 +945,7 @@ @@
             new_name = self.map_tensor_name(merged_name)
             logger.info(f"Repacked {new_name} with shape [{len(experts)}, {shape[0]}, {shape[1]}] and quantization NVFP4")
             self.gguf_writer.add_tensor(new_name, merged, raw_dtype=gguf.GGMLQuantizationType.NVFP4)
+            self._tag_prec_a4(merged_name, new_name)
             scales.sort(key=lambda x: x[0])
             self._write_scales_tensor(new_name.replace(".weight", ".scale"), [s[1] for s in scales])
@@ Expand Down Expand Up / @@ -971,6 +988,9 @@ @@
                 and bool(quant_groups)
                 and all(g.get("format") == "nvfp4-pack-quantized" for g in quant_groups.values() if isinstance(g, dict))
             )
+            self._nvfp4_global_algo = quant_algo
             if quant_algo != "NVFP4":
                 if nvfp4_compressed_tensors:
                     quant_algo = "NVFP4"
@@ Expand All / @@ -980,6 +1000,22 @@ @@
             self._is_nvfp4 = quant_algo in ("NVFP4", "W4A16_NVFP4")
             self._is_mxfp4 = quant_method == "mxfp4"
+            # Per-tensor NVFP4 precision.
+            self._nvfp4_layer_algo = {}
+            if quant_layers:
+                # store all possible module paths and assert if a quantized layer is not in the model
+                modules: set[str] = set()
+                for name in self.model_tensors:
+                    while name := name.rpartition(".")[0]:
+                        modules.add(name)
+                for layer_name, entry in quant_layers.items():
+                    if not isinstance(entry, dict):
+                        continue
+                    if titem := self.filter_tensors((layer_name, lambda: torch.empty(0))):
+                        assert titem[0] in modules, f"quantized_layers entry {layer_name!r} is not in the model tensors"
+                        self._nvfp4_layer_algo[titem[0]] = entry.get("quant_algo")
             # NVFP4 weights are repacked and written directly to gguf_writer.
             # This must run before dequant_model so NVFP4 tensors are removed
             # from model_tensors, leaving only non-NVFP4 (e.g. FP8) for dequant.
@@ Expand Down Expand Up / @@ -1185,6 +1221,12 @@ @@
             logger.info("Set model quantization version")
             self.gguf_writer.add_quantization_version(gguf.GGML_QUANT_VERSION)
+            if self._prec_a4:
+                names = sorted(self._prec_a4.keys())
+                values = [self._prec_a4[n] for n in names]
+                logger.info(f"Set prec_a4 metadata for {len(names)} tensor(s)")
+                self.gguf_writer.add_tensor_extra_prec_a4(names, values)
         def write_vocab(self):
             raise NotImplementedError("write_vocab() must be implemented in subclasses")
@@ Expand Down Expand Up / @@ -1542,15 +1584,15 @@ @@
             from transformers import AutoTokenizer
             tokenizer = AutoTokenizer.from_pretrained(self.dir_model)
             vocab_size = self.hparams.get("vocab_size", len(tokenizer.vocab))  # ty: ignore[unresolved-attribute]
             assert max(tokenizer.vocab.values()) < vocab_size  # ty: ignore[unresolved-attribute]
             tokpre = self.get_vocab_base_pre(tokenizer)
             reverse_vocab = {id_: encoded_tok for encoded_tok, id_ in tokenizer.vocab.items()}  # ty: ignore[unresolved-attribute]
             added_vocab = tokenizer.get_added_vocab()  # ty: ignore[unresolved-attribute]
             added_tokens_decoder = tokenizer.added_tokens_decoder  # ty: ignore[unresolved-attribute]
             for i in range(vocab_size):
                 if i not in reverse_vocab:
@@ Expand All / @@ -1563,7 +1605,7 @@ @@
                         # To avoid unexpected issues - we make sure to normalize non-normalized tokens
                         if not added_tokens_decoder[i].normalized:
                             previous_token = token
                             token = tokenizer.decode(tokenizer.encode(token, add_special_tokens=False))  # ty: ignore[unresolved-attribute, invalid-assignment]
                             if previous_token != token:
                                 logger.info(f"{repr(previous_token)} is encoded and decoded back to {repr(token)} using AutoTokenizer")
@@ Expand Down Expand Up / @@ -1942,10 +1984,10 @@ @@
         def _set_vocab_hybriddna(self):
             from transformers import AutoTokenizer
             tokenizer = AutoTokenizer.from_pretrained(self.dir_model, trust_remote_code=True)
             vocab_size = self.hparams.get("vocab_size", len(tokenizer.vocab))  # ty: ignore[unresolved-attribute]
             assert max(tokenizer.vocab.values()) < vocab_size  # ty: ignore[unresolved-attribute]
             reverse_vocab = {id_: encoded_tok for encoded_tok, id_ in tokenizer.vocab.items()}  # ty: ignore[unresolved-attribute]
             # k-mers can share text with a base-vocab BPE token (e.g. CCCCCC) and get
             # dropped by get_vocab(); a reserved marker suffix (U+E000) keeps each
             # k-mer's own id (llama.cpp strips it on detokenization)
@@ Expand Down @@

docs/build.md

-Original file line number
+Diff line change
@@ Expand Up @@
     Override default, speed-optimized compute types for cuBLAS matrix multiplications.
     Legal values: `auto`, `f16`, `fp16`, `bf16`, `f32`, `fp32`.
+    #### GGML_CUDA_MMQ_PREC
+    Override the activation precision that the model requests for NVFP4 and MXFP4 matrix multiplications.
+    Currently supported values: `auto`, `q8`, `q4`.
+    NVFP4 and MXFP4 layers marked as W4A16 request 8-bit activations, so on Blackwell those layers run through the W4A8 path instead of the native W4A4 path. Set `q4` to keep the native W4A4 path for faster prompt processing at the cost of accuracy, or `q8` to use the W4A8 path for every layer, `auto` uses per-tensor prec metadata (this is the same behavior as when the environment variable is not set).
     ### Unified Memory
     The environment variable `GGML_CUDA_ENABLE_UNIFIED_MEMORY=1` can be used to enable unified memory in Linux. This allows swapping to system RAM instead of crashing when the GPU VRAM is exhausted. In Windows this setting is available in the NVIDIA control panel as `System Memory Fallback`.
@@ Expand Down @@

ggml/src/ggml-cuda/mmq-load-tiles.cuh

            
                      Original file line number
                      Diff line number
                      Diff line change
                  
    @@ -1631,16 +1631,16 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
  
    template <ggml_type type, int J, bool fallback> static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_mxfp4_fp4(

            const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) {

        constexpr int warp_size   = ggml_cuda_get_physical_warp_size();

        constexpr int nwarps      = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size;

        constexpr int I           = ggml_cuda_mmq_get_I(type, J, fallback);

        constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback);

        constexpr int nwarps      = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q4) / warp_size;

        constexpr int I           = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q4);

        constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback, GGML_PREC_Q4);

        int *      x_qs = (int *) x_tile;

        uint32_t * x_sc = (uint32_t *) (x_qs + 2 * MMQ_TILE_NE_K);

        const int txi = threadIdx.x;

        constexpr int iter_k = ggml_cuda_mmq_get_K_vram(type, J, fallback);

        constexpr int iter_k = ggml_cuda_mmq_get_K_vram(type, J, fallback, GGML_PREC_Q4);

        constexpr int threads_per_row = iter_k / QK_MXFP4;  // each thread processes 1 block

        constexpr int rows_per_warp   = warp_size / threads_per_row;

    @@ -1670,12 +1670,12 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
  
        }

    }

    template <ggml_type type, int J, bool fallback> static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_nvfp4(

    template <ggml_type type, int J, bool fallback, ggml_prec prec_src1 = GGML_PREC_Q8> static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_nvfp4(

            const char * __restrict__ x, int * __restrict__ x_tile, const int kb0, const int i_max, const int stride) {

        constexpr int warp_size   = ggml_cuda_get_physical_warp_size();

        constexpr int nwarps      = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size;

        constexpr int I           = ggml_cuda_mmq_get_I(type, J, fallback);

        constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback);

        constexpr int nwarps      = ggml_cuda_mmq_get_nthreads(type, J, fallback, prec_src1) / warp_size;

        constexpr int I           = ggml_cuda_mmq_get_I(type, J, fallback, prec_src1);

        constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback, prec_src1);

    #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)

        int   * x_qs = (int   *) x_tile;

    @@ -1729,12 +1729,12 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
  
    template <ggml_type type, int J, bool fallback> static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_nvfp4_nvfp4(

            const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) {

        constexpr int warp_size       = ggml_cuda_get_physical_warp_size();

        constexpr int nwarps          = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size;

        constexpr int I               = ggml_cuda_mmq_get_I(type, J, fallback);

        constexpr int iter_k          = ggml_cuda_mmq_get_K_vram(type, J, fallback);

        constexpr int nwarps          = ggml_cuda_mmq_get_nthreads(type, J, fallback, GGML_PREC_Q4) / warp_size;

        constexpr int I               = ggml_cuda_mmq_get_I(type, J, fallback, GGML_PREC_Q4);

        constexpr int iter_k          = ggml_cuda_mmq_get_K_vram(type, J, fallback, GGML_PREC_Q4);

        constexpr int threads_per_row = iter_k / QK_NVFP4; // each thread processes 1 block

        constexpr int rows_per_warp   = warp_size / threads_per_row;

        constexpr int sram_stride     = ggml_cuda_mmq_get_sram_stride(type, J, fallback);

        constexpr int sram_stride     = ggml_cuda_mmq_get_sram_stride(type, J, fallback, GGML_PREC_Q4);

        uint32_t * x_u32 = (uint32_t *) x_tile;

ggml/src/ggml-cuda/mmq-vec-dot.cuh

-Original file line number
+Diff line change
@@ Expand Up @@
     }
     // Used for Q3_K, IQ2_S, and IQ2_XS:
-    template <ggml_type type, int J, bool fallback> static __device__ __forceinline__ void ggml_cuda_mmq_vec_dot_q8_0_16_q8_1_mma(
+    template <ggml_type type, int J, bool fallback, ggml_prec prec_src1 = GGML_PREC_Q8> static __device__ __forceinline__ void ggml_cuda_mmq_vec_dot_q8_0_16_q8_1_mma(
             const int * __restrict__ x, const int * __restrict__ y, float * __restrict__ sum, const int k00) {
     #if defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
         constexpr data_layout input_layout = get_input_data_layout();
         typedef tile<16,  4, int, input_layout>        tile_A;
         typedef tile<16,  4, int, input_layout>        tile_B;
         typedef tile<16, 16, int, DATA_LAYOUT_J_MAJOR> tile_C;
-        constexpr int sram_stride   = ggml_cuda_mmq_get_sram_stride(type, J, fallback);
+        constexpr int sram_stride   = ggml_cuda_mmq_get_sram_stride(type, J, fallback, prec_src1);
         constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback);
         constexpr int ntx           = rows_per_warp/tile_C::I; // Number of x minitiles per warp.
@@ Expand Down Expand Up @@
         typedef tile< 8, 4, int> tile_B;
         typedef tile<16, 8, int> tile_C;
-        constexpr int sram_stride   = ggml_cuda_mmq_get_sram_stride(type, J, fallback);
+        constexpr int sram_stride   = ggml_cuda_mmq_get_sram_stride(type, J, fallback, prec_src1);
         constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback);
         constexpr int ntx           = rows_per_warp/tile_C::I; // Number of x minitiles per warp.
@@ Expand Down Expand Up @@
         typedef tile<8,  8, int>   tile_B;
         typedef tile<16, 8, float> tile_C;
-        constexpr int sram_stride   = ggml_cuda_mmq_get_sram_stride(type, J, fallback);
+        constexpr int sram_stride   = ggml_cuda_mmq_get_sram_stride(type, J, fallback, GGML_PREC_Q4);
         constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback);
         constexpr int ntx           = rows_per_warp / tile_C::I;
         constexpr int nfrags        = MMQ_TILE_NE_K / tile_A::J;
@@ Expand Down @@

ggml/src/ggml-cuda/mmq.cu

            
                      Original file line number
                      Diff line number
                      Diff line change
                  
    @@ -5,7 +5,7 @@
  
    #include <cstdint>

    static void ggml_cuda_mul_mat_q_switch_type(ggml_backend_cuda_context & ctx, const mmq_args & args, cudaStream_t stream) {

    static void ggml_cuda_mul_mat_q_switch_type(ggml_backend_cuda_context & ctx, const mmq_args & args, cudaStream_t stream, const ggml_prec prec_src1) {

        switch (args.type_x) {

            case GGML_TYPE_Q1_0:

                mul_mat_q_case<GGML_TYPE_Q1_0>(ctx, args, stream);

    @@ -71,9 +71,18 @@ static void ggml_cuda_mul_mat_q_switch_type(ggml_backend_cuda_context & ctx, con
  
                break;

    // -----------------------------------------------------------------------

            case GGML_TYPE_MXFP4:

                // src1 at Q4 uses the native FP4 instructions, which are Blackwell-only

                if (prec_src1 == GGML_PREC_Q4) {

                    mul_mat_q_case<GGML_TYPE_MXFP4, GGML_PREC_Q4>(ctx, args, stream);

                    break;

                }

                mul_mat_q_case<GGML_TYPE_MXFP4>(ctx, args, stream);

                break;

            case GGML_TYPE_NVFP4:

                if (prec_src1 == GGML_PREC_Q4) {

                    mul_mat_q_case<GGML_TYPE_NVFP4, GGML_PREC_Q4>(ctx, args, stream);

                    break;

                }

                mul_mat_q_case<GGML_TYPE_NVFP4>(ctx, args, stream);

                break;

            default:

    @@ -82,6 +91,47 @@ static void ggml_cuda_mul_mat_q_switch_type(ggml_backend_cuda_context & ctx, con
  
        }

    }

    // overrides the src1 precision requested by the graph, "auto" keeps the requested one

    static ggml_prec ggml_cuda_mmq_get_prec_env() {

        const char * env_c = getenv("GGML_CUDA_MMQ_PREC");

        if (env_c == nullptr) {

            return GGML_PREC_UNDEFINED;

        }

        std::string env_cpp = env_c;

        for (char & c : env_cpp) {

            c = std::tolower(c);

        }

        if (env_cpp == "q4") {

            return GGML_PREC_Q4;

        }

        if (env_cpp == "q8") {

            return GGML_PREC_Q8;

        }

        if (env_cpp != "auto") {

            GGML_LOG_WARN("%s: Unknown value for GGML_CUDA_MMQ_PREC: '%s'. Available: 'q4', 'q8', 'auto'.\n", __func__, env_cpp.c_str());

        }

        return GGML_PREC_UNDEFINED;

    }

    // src1 is quantized to Q8_1 unless the FP4 types can use 4-bit activations, in which case they

    // default to the native W4A4 instructions on Blackwell.

    static ggml_prec ggml_cuda_mmq_get_prec_src1(const ggml_tensor * src0, const ggml_tensor * dst, const int cc) {

        static const ggml_prec prec_env = ggml_cuda_mmq_get_prec_env();

        ggml_prec prec = prec_env;

        if (prec == GGML_PREC_UNDEFINED) {

            prec = (ggml_prec) ggml_get_op_params_i32(dst, 3);

        }

        // Q4 only for the FP4 types on Blackwell

        GGML_ASSERT(prec == GGML_PREC_UNDEFINED || prec == GGML_PREC_Q8 || prec == GGML_PREC_Q4);

        const bool can_use_q4 = (src0->type == GGML_TYPE_NVFP4 || src0->type == GGML_TYPE_MXFP4) && blackwell_mma_available(cc);

        if (prec == GGML_PREC_Q8 || !can_use_q4) {

            return GGML_PREC_Q8;

        }

        return GGML_PREC_Q4;

    }

    void ggml_cuda_mul_mat_q(

            ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, const ggml_tensor * ids, ggml_tensor * dst) {

        GGML_ASSERT(        src1->type == GGML_TYPE_F32);

    @@ -128,7 +178,9 @@ void ggml_cuda_mul_mat_q(
  
        const bool fallback = ne01 % 128 != 0;

        const bool use_native_fp4 = blackwell_mma_available(cc) && (src0->type == GGML_TYPE_MXFP4 || src0->type == GGML_TYPE_NVFP4);

        const ggml_prec prec_src1 = ggml_cuda_mmq_get_prec_src1(src0, dst, cc);

        const bool use_native_fp4 = prec_src1 == GGML_PREC_Q4;

        const size_t y_block_size       = use_native_fp4 ? sizeof(block_fp4_mmq) : sizeof(block_q8_1_mmq);

        const size_t y_values_per_block = use_native_fp4 ? QK_FP4_MMQ            : QK8_1_MMQ;

    @@ -172,7 +224,7 @@ void ggml_cuda_mul_mat_q(
  
                ne02, ne12, s02, s12, s2,

                ne03, ne13, s03, s13, s3,

                ne1, ne1};

            ggml_cuda_mul_mat_q_switch_type(ctx, args, stream);

            ggml_cuda_mul_mat_q_switch_type(ctx, args, stream, prec_src1);

            return;

        }

    @@ -260,7 +312,7 @@ void ggml_cuda_mul_mat_q(
  
            ne03, ne13, s03, s13, s3,

            ne12, ncols_opt};

        ggml_cuda_mul_mat_q_switch_type(ctx, args, stream);

        ggml_cuda_mul_mat_q_switch_type(ctx, args, stream, prec_src1);

    }

    bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t n_experts) {

Navigation Menu

Force NVFP4 W4A8 path for NVFP4_W4A16 layers on Blackwell, where NVFP4 normally uses the native W4A4 path. #24364

Diff view

Diff view

There are no files selected for viewing

GitHub Actions / python type-check

GitHub Actions / python type-check

GitHub Actions / python type-check

GitHub Actions / python type-check

GitHub Actions / python type-check

GitHub Actions / python type-check

GitHub Actions / python type-check

GitHub Actions / python type-check

GitHub Actions / python type-check

GitHub Actions / python type-check

Uh oh!

Uh oh!

Uh oh!

Uh oh!