Add GGUF Q1_0 kernel support (MMQ/MMVQ + enum patch)

python/sglang/srt/layers/quantization/gguf.py

@@ -11,6 +11,10 @@
 from gguf import GGMLQuantizationType as WeightType
 from torch.nn.parameter import Parameter, UninitializedParameter
 
+from sglang.srt.utils.gguf_compat import GGML_TYPE_Q1_0, ensure_q1_0_gguf_compat
+
+ensure_q1_0_gguf_compat()
+
 from sglang.srt.layers.linear import LinearBase
 from sglang.srt.layers.moe import MoeRunnerConfig
 from sglang.srt.layers.quantization.base_config import (
@@ -134,18 +138,22 @@ def is_layer_skipped_gguf(prefix: str, modules_to_not_convert: list[str]):
     WeightType.IQ4_XS,
     WeightType.IQ4_NL,
 }
+# Q1_0 (1-bit Bonsai) — uses raw type ID since it's not yet in the gguf enum
+Q1_0_TYPES = {
+    GGML_TYPE_Q1_0,
+}
 # TODO(Isotr0py): Currently, we don't have MMQ kernel for I-Matrix quantization.
 # Consolidate DEQUANT_TYPES, MMVQ_QUANT_TYPES and MMQ_QUANT_TYPES after we add
 # MMQ kernel for I-Matrix quantization.
-DEQUANT_TYPES = STANDARD_QUANT_TYPES | KQUANT_TYPES | IMATRIX_QUANT_TYPES
-MMVQ_QUANT_TYPES = STANDARD_QUANT_TYPES | KQUANT_TYPES | IMATRIX_QUANT_TYPES
-MMQ_QUANT_TYPES = STANDARD_QUANT_TYPES | KQUANT_TYPES
+DEQUANT_TYPES = STANDARD_QUANT_TYPES | KQUANT_TYPES | IMATRIX_QUANT_TYPES | Q1_0_TYPES
+MMVQ_QUANT_TYPES = STANDARD_QUANT_TYPES | KQUANT_TYPES | IMATRIX_QUANT_TYPES | Q1_0_TYPES
+MMQ_QUANT_TYPES = STANDARD_QUANT_TYPES | KQUANT_TYPES | Q1_0_TYPES
 
 
 def fused_mul_mat_gguf(
     x: torch.Tensor, qweight: torch.Tensor, qweight_type: int
 ) -> torch.Tensor:
-    if qweight_type in IMATRIX_QUANT_TYPES:
+    if qweight_type in Q1_0_TYPES or qweight_type in IMATRIX_QUANT_TYPES:
         mmvq_safe = 8 if qweight.shape[0] > 5120 else 16
     else:
         mmvq_safe = 2 if qweight.shape[0] > 5120 else 6

python/sglang/srt/model_loader/loader.py

@@ -2022,6 +2022,10 @@ def _get_gguf_weights_map(self, model_config: ModelConfig):
                 "Please install gguf via `pip install gguf` to use gguf quantizer."
             ) from err
 
+        from sglang.srt.utils.gguf_compat import ensure_q1_0_gguf_compat
+
+        ensure_q1_0_gguf_compat()
+
         config = model_config.hf_config
         model_type = config.model_type
         # hack: ggufs have a different name than transformers

python/sglang/srt/model_loader/weight_utils.py

@@ -1078,6 +1078,10 @@ def get_gguf_extra_tensor_names(
 ) -> List[str]:
     import gguf
 
+    from sglang.srt.utils.gguf_compat import ensure_q1_0_gguf_compat
+
+    ensure_q1_0_gguf_compat()
+
     reader = gguf.GGUFReader(gguf_file)
     expected_gguf_keys = set(gguf_to_hf_name_map.keys())
     exact_gguf_keys = set([tensor.name for tensor in reader.tensors])
@@ -1095,6 +1099,10 @@ def gguf_quant_weights_iterator(
 
     import gguf
 
+    from sglang.srt.utils.gguf_compat import ensure_q1_0_gguf_compat
+
+    ensure_q1_0_gguf_compat()
+
     reader = gguf.GGUFReader(gguf_file)
 
     for tensor in reader.tensors:

python/sglang/srt/utils/gguf_compat.py

@@ -0,0 +1,84 @@
+# SPDX-License-Identifier: Apache-2.0
+"""Monkey-patch the gguf Python library to recognize Q1_0 if it's not in the
+released pip package yet. This is a temporary shim — remove once the upstream
+gguf pip package includes Q1_0 natively."""
+
+import logging
+
+import numpy as np
+
+logger = logging.getLogger(__name__)
+
+GGML_TYPE_Q1_0 = 41  # matches ggml.h GGML_TYPE_Q1_0
+Q1_0_BLOCK_SIZE = 128
+Q1_0_TYPE_SIZE = 18  # 2 bytes fp16 scale + 16 bytes packed bits
+
+
+def _dequantize_q1_0_blocks(blocks: np.ndarray) -> np.ndarray:
+    """Dequantize Q1_0 blocks using numpy.
+
+    Each block: 2 bytes fp16 scale, 16 bytes packed 1-bit quants (128 bits).
+    Bit value 1 -> +scale, 0 -> -scale.
+    """
+    n_blocks = blocks.shape[0]
+    scales = blocks[:, :2].view(np.float16).astype(np.float32)  # (n_blocks, 1)
+    packed = blocks[:, 2:]  # (n_blocks, 16)
+    # Unpack bits: each byte -> 8 bits, little-endian bit order
+    signs = np.unpackbits(packed, axis=1, bitorder="little")[
+        :, :Q1_0_BLOCK_SIZE
+    ]  # (n_blocks, 128)
+    # Map: 1 -> +scale, 0 -> -scale
+    result = np.where(signs == 1, scales, -scales)
+    return result.astype(np.float32)
+
+
+def ensure_q1_0_gguf_compat():
+    """Register Q1_0 type in the gguf library if not already present."""
+    import gguf
+    import gguf.quants as gguf_quants
+
+    if getattr(gguf, "_sglang_q1_0_compat", False):
+        return
+
+    # Check if Q1_0 is already natively supported
+    from gguf import GGMLQuantizationType
+
+    if hasattr(GGMLQuantizationType, "Q1_0"):
+        logger.debug("Q1_0 already in gguf library, skipping compat shim")
+        gguf._sglang_q1_0_compat = True
+        return
+
+    logger.info("Applying Q1_0 compat shim to gguf library")
+
+    # 0. Add Q1_0 to the GGMLQuantizationType enum so GGUFReader can parse it
+    from gguf import GGMLQuantizationType
+    import enum
+
+    # Extend the enum with Q1_0 = 41
+    new_member = int.__new__(GGMLQuantizationType, GGML_TYPE_Q1_0)
+    new_member._name_ = "Q1_0"
+    new_member._value_ = GGML_TYPE_Q1_0
+    GGMLQuantizationType._member_map_["Q1_0"] = new_member
+    GGMLQuantizationType._value2member_map_[GGML_TYPE_Q1_0] = new_member
+    GGMLQuantizationType._member_names_.append("Q1_0")
+
+    # 1. Register quant sizes (block_size, type_size)
+    gguf.GGML_QUANT_SIZES[GGML_TYPE_Q1_0] = (Q1_0_BLOCK_SIZE, Q1_0_TYPE_SIZE)
+    gguf_quants.GGML_QUANT_SIZES[GGML_TYPE_Q1_0] = (
+        Q1_0_BLOCK_SIZE,
+        Q1_0_TYPE_SIZE,
+    )
+
+    # 2. Create a quant class for dequantization support
+    class Q1_0(gguf_quants.__Quant, qtype=GGML_TYPE_Q1_0):
+        block_size = Q1_0_BLOCK_SIZE
+        type_size = Q1_0_TYPE_SIZE
+
+        @classmethod
+        def dequantize_blocks(cls, blocks: np.ndarray) -> np.ndarray:
+            return _dequantize_q1_0_blocks(blocks)
+
+    gguf_quants._type_traits[GGML_TYPE_Q1_0] = Q1_0
+
+    # 3. Mark as applied
+    gguf._sglang_q1_0_compat = True

python/sglang/srt/utils/hf_transformers/config.py

@@ -59,6 +59,11 @@ def get_config(
     is_gguf = check_gguf_file(model)
     if is_gguf:
         _ensure_gguf_version()
+        # Ensure Q1_0 gguf compat is applied before transformers reads the GGUF
+        if str(model).endswith(".gguf"):
+            from sglang.srt.utils.gguf_compat import ensure_q1_0_gguf_compat
+
+            ensure_q1_0_gguf_compat()
         kwargs["gguf_file"] = model
         model = Path(model).parent
 

sgl-kernel/csrc/quantization/gguf/dequantize.cuh

@@ -2,6 +2,25 @@
 // https://github.com/vllm-project/vllm/blob/4492e3a55428e161ca8db381edc28263e5da4c8d/csrc/quantization/gguf/dequantize.cuh
 // copied and adapted from https://github.com/ggerganov/llama.cpp/blob/b2899/ggml-cuda/convert.cu
 // Dequant functions
+static __device__ __forceinline__ void dequantize_q1_0(const void* vx, const int ib, const int iqs, dfloat2& v) {
+  const block_q1_0* x = (const block_q1_0*)vx;
+
+  const dfloat d = x[ib].d;
+
+  const int byte_index_0 = iqs / 8;
+  const int bit_offset_0 = iqs % 8;
+  const int byte_index_1 = (iqs + 1) / 8;
+  const int bit_offset_1 = (iqs + 1) % 8;
+
+  // Extract bits: 1 = +d, 0 = -d (branchless)
+  const int bit_0 = (x[ib].qs[byte_index_0] >> bit_offset_0) & 1;
+  const int bit_1 = (x[ib].qs[byte_index_1] >> bit_offset_1) & 1;
+
+  v.x = __int2half_rn(2 * bit_0 - 1);
+  v.y = __int2half_rn(2 * bit_1 - 1);
+  v = __hmul2(v, {d, d});
+}
+
 static __device__ __forceinline__ void dequantize_q4_0(const void* vx, const int ib, const int iqs, dfloat2& v) {
   const block_q4_0* x = (const block_q4_0*)vx;
 
@@ -577,6 +596,8 @@ static to_cuda_ggml_t<dst_t> ggml_get_to_cuda(int64_t type) {
       return dequantize_row_iq4_xs_cuda;
     case 29:
       return dequantize_row_iq1_m_cuda;
+    case 41:
+      return dequantize_block_cuda<QK1_0, QR1_0, dequantize_q1_0>;
     default:
       return nullptr;
   }

sgl-kernel/csrc/quantization/gguf/ggml-common.h

@@ -15,6 +15,19 @@
 // QR = QK / number of values before dequantization
 // QI = number of 32 bit integers before dequantization
 
+// Q1_0: 1-bit quantization, 128 elements per block
+// Block layout: [half d (2 bytes)] [uint8_t qs[16] (16 bytes)] = 18 bytes total
+// Each bit in qs: 1 -> +d, 0 -> -d (symmetric binary quantization)
+#define QK1_0 128
+#define QR1_0 1
+#define QI1_0 4  // 128 elements / 32 per Q8_1 block = 4 Q8_1 blocks per Q1_0 block
+typedef struct {
+  half d;                  // scale factor
+  uint8_t qs[QK1_0 / 8];  // 1-bit quants (16 bytes = 128 bits)
+} block_q1_0;
+static_assert(sizeof(block_q1_0) == sizeof(half) + QK1_0 / 8,
+              "wrong q1_0 block size/padding");
+
 #define QK4_0 32
 #define QR4_0 2
 #define QI4_0 (QK4_0 / (4 * QR4_0))

sgl-kernel/csrc/quantization/gguf/gguf_kernel.cu

@@ -184,6 +184,10 @@ torch::Tensor ggml_mul_mat_vec_a8(
         mul_mat_vec_iq1_m_q8_1_cuda<scalar_t>(
             (void*)W.data_ptr(), (void*)quant_X.data_ptr(), (scalar_t*)Y.data_ptr(), col, row, vecs, stream);
         break;
+      case 41:
+        mul_mat_vec_q1_0_q8_1_cuda<scalar_t>(
+            (void*)W.data_ptr(), (void*)quant_X.data_ptr(), (scalar_t*)Y.data_ptr(), col, row, vecs, stream);
+        break;
     }
   });
   return Y;
@@ -327,6 +331,18 @@ torch::Tensor ggml_mul_mat_a8(
             row,
             stream);
         break;
+      case 41:
+        ggml_mul_mat_q1_0_q8_1_cuda(
+            (void*)W.data_ptr(),
+            (void*)quant_X.data_ptr(),
+            (scalar_t*)Y.data_ptr(),
+            col,
+            row,
+            batch,
+            padded,
+            row,
+            stream);
+        break;
     }
   });
   return Y;
@@ -804,6 +820,19 @@ torch::Tensor ggml_moe_a8_vec(
             quant_X.stride(0),
             stream);
         break;
+      case 41:
+        moe_vec_q1_0_q8_1_cuda<scalar_t>(
+            (void*)W.data_ptr(),
+            (void*)quant_X.data_ptr(),
+            (scalar_t*)Y.data_ptr(),
+            (int*)topk_ids.data_ptr(),
+            top_k,
+            tokens,
+            col,
+            row,
+            quant_X.stride(0),
+            stream);
+        break;
     }
   });
   return Y;

sgl-kernel/csrc/quantization/gguf/mmq.cuh

@@ -504,6 +504,84 @@ static void ggml_mul_mat_q8_0_q8_1_cuda(
   }
 }
 
+// Q1_0 MMQ — uses Q8_0 tile layout and dot product, with custom load_tiles
+#if defined(USE_ROCM)
+#define MMQ_X_Q1_0 64
+#define MMQ_Y_Q1_0 128
+#define NWARPS_Q1_0 8
+#else
+#define MMQ_X_Q1_0 4
+#define MMQ_Y_Q1_0 32
+#define NWARPS_Q1_0 4
+#endif
+
+template <typename scalar_t, bool need_check>
+static __global__ void
+#if defined(USE_ROCM)
+__launch_bounds__(WARP_SIZE_GGUF* NWARPS_Q1_0, 2)
+#endif
+    mul_mat_q1_0(
+        const void* __restrict__ vx,
+        const void* __restrict__ vy,
+        scalar_t* __restrict__ dst,
+        const int ncols_x,
+        const int nrows_x,
+        const int ncols_y,
+        const int nrows_y,
+        const int nrows_dst) {
+  const int mmq_x = MMQ_X_Q1_0;
+  const int mmq_y = MMQ_Y_Q1_0;
+  const int nwarps = NWARPS_Q1_0;
+
+  // QI1_0_MMQ=32: blocks_per_warp=1, each iteration = 1 Q1_0 block = 128 elements
+  // Reuses Q8_0 vec_dot since unpacked data is in the same signed-byte format
+  mul_mat_q<
+      scalar_t,
+      QK1_0,
+      QR1_0,
+      QI1_0_MMQ,
+      false,
+      block_q1_0,
+      mmq_x,
+      mmq_y,
+      nwarps,
+      allocate_tiles_q1_0<mmq_y>,
+      load_tiles_q1_0<mmq_y, nwarps, need_check>,
+      VDR_Q8_0_Q8_1_MMQ,
+      vec_dot_q8_0_q8_1_mul_mat>(vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, nrows_dst);
+}
+
+template <typename scalar_t>
+static void ggml_mul_mat_q1_0_q8_1_cuda(
+    const void* vx,
+    const void* vy,
+    scalar_t* dst,
+    const int ncols_x,
+    const int nrows_x,
+    const int ncols_y,
+    const int nrows_y,
+    const int nrows_dst,
+    cudaStream_t stream) {
+  const int mmq_x = MMQ_X_Q1_0;
+  const int mmq_y = MMQ_Y_Q1_0;
+  const int nwarps = NWARPS_Q1_0;
+
+  const int block_num_x = (nrows_x + mmq_y - 1) / mmq_y;
+  const int block_num_y = (ncols_y + mmq_x - 1) / mmq_x;
+  const dim3 block_nums(block_num_x, block_num_y, 1);
+  const dim3 block_dims(WARP_SIZE_GGUF, nwarps, 1);
+
+  if (nrows_x % mmq_y == 0) {
+    const bool need_check = false;
+    mul_mat_q1_0<scalar_t, need_check>
+        <<<block_nums, block_dims, 0, stream>>>(vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, nrows_dst);
+  } else {
+    const bool need_check = true;
+    mul_mat_q1_0<scalar_t, need_check>
+        <<<block_nums, block_dims, 0, stream>>>(vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, nrows_dst);
+  }
+}
+
 #if defined(USE_ROCM)
 #define MMQ_X_Q2_K 64
 #define MMQ_Y_Q2_K 128

sgl-kernel/csrc/quantization/gguf/mmvq.cuh

@@ -47,6 +47,22 @@ static __global__ void mul_mat_vec_q(
   }
 }
 
+template <typename scalar_t>
+static void mul_mat_vec_q1_0_q8_1_cuda(
+    const void* vx,
+    const void* vy,
+    scalar_t* dst,
+    const int ncols,
+    const int nrows,
+    const int nvecs,
+    cudaStream_t stream) {
+  const int block_num_y = (nrows + GGML_CUDA_MMV_Y - 1) / GGML_CUDA_MMV_Y;
+  const dim3 block_nums(block_num_y, nvecs, 1);
+  const dim3 block_dims(WARP_SIZE, GGML_CUDA_MMV_Y, 1);
+  mul_mat_vec_q<scalar_t, QK1_0, QI1_0, block_q1_0, VDR_Q1_0_Q8_1_MMVQ, vec_dot_q1_0_q8_1>
+      <<<block_nums, block_dims, 0, stream>>>(vx, vy, dst, ncols, nrows, nvecs);
+}
+
 template <typename scalar_t>
 static void mul_mat_vec_q4_0_q8_1_cuda(
     const void* vx,