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

conversion/base.py

@@ -160,6 +160,9 @@ def __init__(self, dir_model: Path, ftype: gguf.LlamaFileType, fname_out: Path,
         self.dir_model_card = dir_model  # overridden in convert_lora_to_gguf.py
         self._is_nvfp4 = False
         self._is_mxfp4 = False
+        self._nvfp4_all_w4a16 = False # checkpoint quantized entirely as W4A16_NVFP4
+        self._prec_a8_w4a4_override: set[str] = set() # gguf names explicitly kept at plain NVFP4 (W4A4) under a global W4A16 default
+        self._allow_prec_a8: dict[str, bool] = {} # gguf tensor name -> wants = 8-bit (A8) activations
         self._fp8_as_q8 = fp8_as_q8
         self._fp8_dequantized: set[str] = set()
 
@@ -617,6 +620,47 @@ def map_tensor_name(self, name: str, try_suffixes: Sequence[str] = (".weight", "
             raise ValueError(f"Can not map tensor {name!r}")
         return new_name
 
+    def _gguf_weight_name(self, name: str) -> str:
+        if name.endswith((".weight", ".bias")):
+            return name
+        return name + ".weight"
+
+    def _tag_prec_a8(self, new_name: str) -> None:
+        # Tag a W4A16_NVFP4 weight as needing 8-bit activation.
+        if self._nvfp4_all_w4a16:
+            gguf_name = self._gguf_weight_name(new_name)
+            if gguf_name not in self._prec_a8_w4a4_override:
+                self._allow_prec_a8[gguf_name] = True
+
+    def _hf_quant_tensors_to_gguf(self, hf_name: str) -> list[str]:
+        # Map an HF quantized-layer name to its GGUF tensor name(s).
+        if hf_name == "lm_head" or hf_name.endswith(".lm_head"):
+            return ["output.weight"]
+
+        name = hf_name
+        if name.startswith("model.language_model."):
+            name = "model." + name[len("model.language_model."):]
+
+        m = re.fullmatch(r"model\.layers\.(\d+)\.mlp\.experts", name)
+        if m:
+            bid = int(m.group(1))
+            return [
+                self._gguf_weight_name(self.format_tensor_name(gguf.MODEL_TENSOR.FFN_GATE_EXP, bid)),
+                self._gguf_weight_name(self.format_tensor_name(gguf.MODEL_TENSOR.FFN_UP_EXP, bid)),
+                self._gguf_weight_name(self.format_tensor_name(gguf.MODEL_TENSOR.FFN_DOWN_EXP, bid)),
+            ]
+
+        candidates = [name]
+        if not name.endswith((".weight", ".bias")):
+            candidates.append(name + ".weight")
+
+        for cand in candidates:
+            try:
+                return [self._gguf_weight_name(self.map_tensor_name(cand))]
+            except ValueError:
+                continue
+        return []
+
     def set_gguf_parameters(self):
         raise NotImplementedError("set_gguf_parameters() must be implemented in subclasses")
 
@@ -726,6 +770,7 @@ def _repack_nvfp4(self, name: str, weight: Tensor, scale: Tensor, scale2: Tensor
         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_a8(new_name)
 
         self._write_scale_tensor(new_name.replace(".weight", ".scale"), scale2)
         self._write_scale_tensor(new_name.replace(".weight", ".input_scale"), input_scale)
@@ -818,6 +863,7 @@ def _flush_nvfp4_experts(self, key, expert_blocks, expert_scales, expert_input_s
         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_a8(new_name)
 
         scales.sort(key=lambda x: x[0])
         self._write_scales_tensor(new_name.replace(".weight", ".scale"), [s[1] for s in scales])
@@ -860,6 +906,9 @@ def prepare_tensors(self):
             and bool(quant_groups)
             and all(g.get("format") == "nvfp4-pack-quantized" for g in quant_groups.values() if isinstance(g, dict))
         )
+        # A checkpoint quantized entirely as W4A16_NVFP4 uses a single global quant_algo
+        self._nvfp4_all_w4a16 = quant_algo == "W4A16_NVFP4"
+
         if quant_algo != "NVFP4":
             if nvfp4_compressed_tensors:
                 quant_algo = "NVFP4"
@@ -869,6 +918,21 @@ def prepare_tensors(self):
         self._is_nvfp4 = quant_algo in ("NVFP4", "W4A16_NVFP4")
         self._is_mxfp4 = quant_method == "mxfp4"
 
+        # Collect per-tensor W4A16_NVFP4 metadata (activations were not quantized to 4-bit).
+        if self._is_nvfp4 and not self._nvfp4_all_w4a16:
+            # Per-layer map tag only the W4A16_NVFP4 layers.
+            for tensor_name, entry in quant_layers.items():
+                if not isinstance(entry, dict) or entry.get("quant_algo") != "W4A16_NVFP4":
+                    continue
+                for gguf_name in self._hf_quant_tensors_to_gguf(tensor_name):
+                    self._allow_prec_a8[gguf_name] = True
+        elif self._nvfp4_all_w4a16:
+            # Record any per-layer entries that override back to plain NVFP4 (W4A4)
+            for tensor_name, entry in quant_layers.items():
+                algo = entry.get("quant_algo") if isinstance(entry, dict) else None
+                if isinstance(algo, str) and algo.endswith("NVFP4") and algo != "W4A16_NVFP4":
+                    self._prec_a8_w4a4_override.update(self._hf_quant_tensors_to_gguf(tensor_name))
+
         # 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.
@@ -1061,6 +1125,12 @@ def prepare_metadata(self, vocab_only: bool):
         logger.info("Set model quantization version")
         self.gguf_writer.add_quantization_version(gguf.GGML_QUANT_VERSION)
 
+        if self._allow_prec_a8:
+            names = sorted(self._allow_prec_a8.keys())
+            values = [self._allow_prec_a8[n] for n in names]
+            logger.info(f"Set allow_prec_a8 metadata for {len(names)} tensor(s)")
+            self.gguf_writer.add_tensor_extra_allow_prec_a8(names, values)
+
     def write_vocab(self):
         raise NotImplementedError("write_vocab() must be implemented in subclasses")
 

docs/build.md

@@ -281,6 +281,10 @@ Consider setting `CUDA_SCALE_LAUNCH_QUEUES=4x`, which increases the CUDA command
 Override default, speed-optimized compute types for cuBLAS matrix multiplications.
 Legal values: `auto`, `f16`, `fp16`, `bf16`, `f32`, `fp32`.
 
+#### GGML_CUDA_FORCE_W4A4
+
+NVFP4 models that carry W4A16 layers request higher-precision activations (W4A8), so on Blackwell those layers run through the W4A8 path instead of the native W4A4 path. Set `GGML_CUDA_FORCE_W4A4=1` to override that request and keep the native W4A4 path for faster prompt processing at the cost of accuracy.
+
 ### 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`.

ggml/include/ggml.h

@@ -434,9 +434,20 @@ extern "C" {
     };
 
     // precision
+    // this enum is used to declare the allowed floating-point types that can be used during the compute of an op
+    // the declared types can be:
+    //  - result accumulation type
+    //  - source tensor data representation type
+    //  - etc.
+    // the precision parameters are stored as ggml_tensor.op_params to the respective ops
     enum ggml_prec {
-        GGML_PREC_DEFAULT =  0, // stored as ggml_tensor.op_params, 0 by default
-        GGML_PREC_F32     = 10,
+        GGML_PREC_UNDEFINED = 0,
+        GGML_PREC_DEFAULT   = 0,  // note: deprecated, use GGML_PREC_UNDEFINED
+        GGML_PREC_F32       = 10,
+        GGML_PREC_BF16      = 15,
+        GGML_PREC_F16       = 20,
+        GGML_PREC_Q8        = 30,
+        GGML_PREC_Q4        = 40,
     };
 
     // op hint
@@ -1422,6 +1433,37 @@ extern "C" {
             struct ggml_tensor  * b,
             float                 eps);
 
+    // set the minimum required accumulator type for the implementation to use during the compute
+    // for example:
+    //  - GGML_PREC_F32  - requires accumulation of the results in F32
+    //  - GGML_PREC_BF16 - can accumulate the results in BF16, F32
+    //  - GGML_PREC_F16  - can accumulate the results in F16, F32
+    //  - GGML_PREC_Q8   - not allowed
+    //  - GGML_PREC_Q4   - not allowed
+    GGML_API void ggml_prec_set_acc(
+            struct ggml_tensor * a,
+            enum ggml_prec       prec);
+
+    // set the smallest rank that the implementation can use to internally convert the src[idx] data to
+    // ranks in decreasing order:
+    //  - GGML_PREC_F32  - GGML_TYPE_F32
+    //  - GGML_PREC_BF16 - GGML_TYPE_BF16
+    //  - GGML_PREC_F16  - GGML_TYPE_F16,
+    //  - GGML_PREC_Q8   - GGML_TYPE_Q8_0, GGML_TYPE_Q8_1, GGML_TYPE_Q8_K, etc.
+    //  - GGML_PREC_Q4   - GGML_TYPE_Q4_0, GGML_TYPE_Q4_1, GGML_TYPE_Q4_K, GGML_TYPE_NVFP4, GGML_TYPE_MXFP4, etc.
+    //
+    // for example:
+    //   - ggml_prec_set_src(a, GGML_PREC_Q8, 1):
+    //     - allows the implementation to quantize F32, BF16, F16 data of src[1] down to GGML_TYPE_Q8_0
+    //     - cannot quantize it down to GGML_TYPE_Q4_0 or GGML_TYPE_NVFP4
+    //   - ggml_prec_set_src(a, GGML_PREC_Q4, 1):
+    //     - allows the implementation to quantize F32, BF16, F16 data of src[1] down to GGML_TYPE_Q8_0 or GGML_TYPE_NVFP4
+    //
+    GGML_API void ggml_prec_set_src(
+            struct ggml_tensor * a,
+            enum ggml_prec       prec,
+            int                  idx);
+
     // A: k columns, n rows => [ne03, ne02, n, k]
     // B: k columns, m rows  (i.e. we transpose it internally) => [ne03 * x, ne02 * y, m, k]
     // result is n columns, m rows => [ne03 * x, ne02 * y, m, n]
@@ -1432,9 +1474,10 @@ extern "C" {
 
     // change the precision of a matrix multiplication
     // set to GGML_PREC_F32 for higher precision (useful for phi-2)
-    GGML_API void ggml_mul_mat_set_prec(
+    GGML_DEPRECATED(GGML_API void ggml_mul_mat_set_prec(
             struct ggml_tensor * a,
-            enum ggml_prec       prec);
+            enum ggml_prec       prec),
+        "use ggml_prec_set_acc() instead");
 
     // change the hint of a matrix multiplication
     GGML_API void ggml_mul_mat_set_hint(
@@ -2439,9 +2482,10 @@ extern "C" {
             float                 max_bias,
             float                 logit_softcap);
 
-    GGML_API void ggml_flash_attn_ext_set_prec(
+    GGML_DEPRECATED(GGML_API void ggml_flash_attn_ext_set_prec(
             struct ggml_tensor * a,
-            enum ggml_prec       prec);
+            enum ggml_prec       prec),
+        "use ggml_prec_set_acc() instead");
 
     GGML_API enum ggml_prec ggml_flash_attn_ext_get_prec(
             const struct ggml_tensor * a);

ggml/src/ggml-cuda/CMakeLists.txt

@@ -131,6 +131,10 @@ if (CUDAToolkit_FOUND)
 
     add_compile_definitions(GGML_CUDA_PEER_MAX_BATCH_SIZE=${GGML_CUDA_PEER_MAX_BATCH_SIZE})
 
+    if (CMAKE_CUDA_ARCHITECTURES MATCHES "(^|;)12[0-9]a(-real|-virtual)?($|;)")
+        add_compile_definitions(GGML_CUDA_HAS_BLACKWELL_TARGET)
+    endif()
+
     if (GGML_CUDA_GRAPHS)
         add_compile_definitions(GGML_CUDA_USE_GRAPHS)
     endif()

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

@@ -1662,12 +1662,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, bool force_w4a8 = false> 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, force_w4a8) / warp_size;
+    constexpr int I           = ggml_cuda_mmq_get_I(type, J, fallback, force_w4a8);
+    constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback, force_w4a8);
 
 #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
     int   * x_qs = (int   *) x_tile;

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

@@ -478,16 +478,16 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
 }
 
 // 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, bool force_w4a8 = false> 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 I             = ggml_cuda_mmq_get_I(type, J, fallback);
-    constexpr int sram_stride   = ggml_cuda_mmq_get_sram_stride(type, J, fallback);
+    constexpr int I             = ggml_cuda_mmq_get_I(type, J, fallback, force_w4a8);
+    constexpr int sram_stride   = ggml_cuda_mmq_get_sram_stride(type, J, fallback, force_w4a8);
     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.
 
@@ -537,8 +537,8 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
     typedef tile< 8, 4, int> tile_B;
     typedef tile<16, 8, int> tile_C;
 
-    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 I             = ggml_cuda_mmq_get_I(type, J, fallback, force_w4a8);
+    constexpr int sram_stride   = ggml_cuda_mmq_get_sram_stride(type, J, fallback, force_w4a8);
     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.
 

ggml/src/ggml-cuda/mmq.cu

@@ -5,7 +5,8 @@
 
 #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,
+                                            [[maybe_unused]] bool force_w4a8 = false) {
     switch (args.type_x) {
         case GGML_TYPE_Q1_0:
             mul_mat_q_case<GGML_TYPE_Q1_0>(ctx, args, stream);
@@ -74,6 +75,13 @@ static void ggml_cuda_mul_mat_q_switch_type(ggml_backend_cuda_context & ctx, con
             mul_mat_q_case<GGML_TYPE_MXFP4>(ctx, args, stream);
             break;
         case GGML_TYPE_NVFP4:
+#ifdef GGML_CUDA_HAS_BLACKWELL_TARGET
+            // W4A16 NVFP4: dispatch the W4A8 instantiation so activations stay at higher precision even on Blackwell.
+            if (force_w4a8) {
+                mul_mat_q_case<GGML_TYPE_NVFP4, true>(ctx, args, stream);
+                break;
+            }
+#endif // GGML_CUDA_HAS_BLACKWELL_TARGET
             mul_mat_q_case<GGML_TYPE_NVFP4>(ctx, args, stream);
             break;
         default:
@@ -82,6 +90,19 @@ static void ggml_cuda_mul_mat_q_switch_type(ggml_backend_cuda_context & ctx, con
     }
 }
 
+// NVFP4 defaults to native W4A4 on Blackwell. A src1 precision of GGML_PREC_Q8 selects the W4A8
+// path, unless GGML_CUDA_FORCE_W4A4 overrides it.
+static inline bool ggml_cuda_mmq_force_w4a8(const ggml_tensor * src0, const ggml_tensor * dst) {
+    static const bool force_w4a4 = []() {
+        const char * env = getenv("GGML_CUDA_FORCE_W4A4");
+        return env != nullptr && std::atoi(env) != 0;
+    }();
+    if (force_w4a4 || src0->type != GGML_TYPE_NVFP4) {
+        return false;
+    }
+    return ggml_get_op_params_i32(dst, 3) == GGML_PREC_Q8;
+}
+
 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 +149,10 @@ 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 bool force_w4a8 = ggml_cuda_mmq_force_w4a8(src0, dst);
+
+    const bool use_native_fp4 = !force_w4a8 && blackwell_mma_available(cc) &&
+                                (src0->type == GGML_TYPE_MXFP4 || src0->type == GGML_TYPE_NVFP4);
     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 +196,7 @@ void ggml_cuda_mul_mat_q(
             ne02, ne12, s02, s12, s2,
             ne03, ne13, s03, s13, s3,
             ne1};
-        ggml_cuda_mul_mat_q_switch_type(ctx, args, stream);
+        ggml_cuda_mul_mat_q_switch_type(ctx, args, stream, force_w4a8);
         return;
     }
 
@@ -253,7 +277,7 @@ void ggml_cuda_mul_mat_q(
         ne03, ne13, s03, s13, s3,
         ne12};
 
-    ggml_cuda_mul_mat_q_switch_type(ctx, args, stream);
+    ggml_cuda_mul_mat_q_switch_type(ctx, args, stream, force_w4a8);
 }
 
 bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t n_experts) {