From 1cfe32a659b01a3013da8498e74ea7509bedb06b Mon Sep 17 00:00:00 2001 From: Pasha Khosravi Date: Fri, 17 Apr 2026 23:50:15 +0000 Subject: [PATCH] Add GGUF Q1_0 kernel support (MMQ/MMVQ + enum patch) --- python/sglang/srt/layers/quantization/gguf.py | 16 ++- python/sglang/srt/model_loader/loader.py | 4 + .../sglang/srt/model_loader/weight_utils.py | 8 ++ python/sglang/srt/utils/gguf_compat.py | 84 ++++++++++++ .../srt/utils/hf_transformers/config.py | 5 + .../csrc/quantization/gguf/dequantize.cuh | 21 +++ .../csrc/quantization/gguf/ggml-common.h | 13 ++ .../csrc/quantization/gguf/gguf_kernel.cu | 29 ++++ sgl-kernel/csrc/quantization/gguf/mmq.cuh | 78 +++++++++++ sgl-kernel/csrc/quantization/gguf/mmvq.cuh | 16 +++ sgl-kernel/csrc/quantization/gguf/moe_vec.cuh | 19 +++ sgl-kernel/csrc/quantization/gguf/vecdotq.cuh | 128 ++++++++++++++++++ 12 files changed, 417 insertions(+), 4 deletions(-) create mode 100644 python/sglang/srt/utils/gguf_compat.py diff --git a/python/sglang/srt/layers/quantization/gguf.py b/python/sglang/srt/layers/quantization/gguf.py index 748b3c29463d..dc88af122143 100644 --- a/python/sglang/srt/layers/quantization/gguf.py +++ b/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 diff --git a/python/sglang/srt/model_loader/loader.py b/python/sglang/srt/model_loader/loader.py index 5cdb3d4fc6ae..e3b551882c9d 100644 --- a/python/sglang/srt/model_loader/loader.py +++ b/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 diff --git a/python/sglang/srt/model_loader/weight_utils.py b/python/sglang/srt/model_loader/weight_utils.py index ddb7a3c365ad..429a67c422ed 100644 --- a/python/sglang/srt/model_loader/weight_utils.py +++ b/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: diff --git a/python/sglang/srt/utils/gguf_compat.py b/python/sglang/srt/utils/gguf_compat.py new file mode 100644 index 000000000000..4208e50ae8d1 --- /dev/null +++ b/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 diff --git a/python/sglang/srt/utils/hf_transformers/config.py b/python/sglang/srt/utils/hf_transformers/config.py index 759dade67afc..a9f056734188 100644 --- a/python/sglang/srt/utils/hf_transformers/config.py +++ b/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 diff --git a/sgl-kernel/csrc/quantization/gguf/dequantize.cuh b/sgl-kernel/csrc/quantization/gguf/dequantize.cuh index 57cb1a27cfb6..fd0383e54492 100644 --- a/sgl-kernel/csrc/quantization/gguf/dequantize.cuh +++ b/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 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; default: return nullptr; } diff --git a/sgl-kernel/csrc/quantization/gguf/ggml-common.h b/sgl-kernel/csrc/quantization/gguf/ggml-common.h index f6fbe57aaf33..9f52f4b3e812 100644 --- a/sgl-kernel/csrc/quantization/gguf/ggml-common.h +++ b/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)) diff --git a/sgl-kernel/csrc/quantization/gguf/gguf_kernel.cu b/sgl-kernel/csrc/quantization/gguf/gguf_kernel.cu index 1a4955f517a0..e04769fd8e84 100644 --- a/sgl-kernel/csrc/quantization/gguf/gguf_kernel.cu +++ b/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( (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( + (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( + (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; diff --git a/sgl-kernel/csrc/quantization/gguf/mmq.cuh b/sgl-kernel/csrc/quantization/gguf/mmq.cuh index 9838c14f0118..d15b51d94d49 100644 --- a/sgl-kernel/csrc/quantization/gguf/mmq.cuh +++ b/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 +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, + load_tiles_q1_0, + 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 +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 + <<>>(vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, nrows_dst); + } else { + const bool need_check = true; + mul_mat_q1_0 + <<>>(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 diff --git a/sgl-kernel/csrc/quantization/gguf/mmvq.cuh b/sgl-kernel/csrc/quantization/gguf/mmvq.cuh index 7331731ace86..5fcc9a8b7662 100644 --- a/sgl-kernel/csrc/quantization/gguf/mmvq.cuh +++ b/sgl-kernel/csrc/quantization/gguf/mmvq.cuh @@ -47,6 +47,22 @@ static __global__ void mul_mat_vec_q( } } +template +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 + <<>>(vx, vy, dst, ncols, nrows, nvecs); +} + template static void mul_mat_vec_q4_0_q8_1_cuda( const void* vx, diff --git a/sgl-kernel/csrc/quantization/gguf/moe_vec.cuh b/sgl-kernel/csrc/quantization/gguf/moe_vec.cuh index 8cef9e080a22..133578090c38 100644 --- a/sgl-kernel/csrc/quantization/gguf/moe_vec.cuh +++ b/sgl-kernel/csrc/quantization/gguf/moe_vec.cuh @@ -411,3 +411,22 @@ static void moe_vec_iq3_s_q8_1_cuda( moe_vec_q <<>>(vx, vy, dst, topk_ids, top_k, ncols, nrows, token_stride); } + +template +static void moe_vec_q1_0_q8_1_cuda( + const void* vx, + const void* vy, + scalar_t* dst, + const int* topk_ids, + const int top_k, + const int tokens, + const int ncols, + const int nrows, + const int token_stride, + 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, 1, tokens * top_k); + const dim3 block_dims(WARP_SIZE, GGML_CUDA_MMV_Y, 1); + moe_vec_q + <<>>(vx, vy, dst, topk_ids, top_k, ncols, nrows, token_stride); +} diff --git a/sgl-kernel/csrc/quantization/gguf/vecdotq.cuh b/sgl-kernel/csrc/quantization/gguf/vecdotq.cuh index 933c2aa51b71..26155f29a15c 100644 --- a/sgl-kernel/csrc/quantization/gguf/vecdotq.cuh +++ b/sgl-kernel/csrc/quantization/gguf/vecdotq.cuh @@ -42,6 +42,9 @@ static __device__ __forceinline__ int get_int_from_uint8_aligned(const uint8_t* // VDR = vec dot ratio, how many contiguous integers each thread processes when the vec dot kernel is called // MMVQ = mul_mat_vec_q, MMQ = mul_mat_q +#define VDR_Q1_0_Q8_1_MMVQ 1 // Process one 32-element chunk at a time +#define VDR_Q1_0_Q8_1_MMQ 4 // Q1_0 has 128 bits (4 ints) per block + #define VDR_Q4_0_Q8_1_MMVQ 2 #define VDR_Q4_0_Q8_1_MMQ 4 @@ -535,6 +538,54 @@ static __device__ __forceinline__ float vec_dot_q6_K_q8_1_impl_mmq( #endif } +static __device__ __forceinline__ float +vec_dot_q1_0_q8_1(const void* __restrict__ vbq, const block_q8_1* __restrict__ bq8_1, const int& iqs) { +#if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 610 || defined USE_ROCM + const block_q1_0* bq1_0 = (const block_q1_0*)vbq; + + // Q1_0: 128 elements with ONE scale + // Q8_1: 32 elements per block with individual scales + // iqs selects which of the 4 chunks of 32 elements to process (0-3) + const float d1 = __half2float(bq1_0->d); + + // Select the Q8_1 block for this chunk + const block_q8_1* bq8_1_chunk = bq8_1 + iqs; + + // Load 32 bits (4 bytes) for this chunk from Q1_0 + const int offset = iqs * 4; + const int v = bq1_0->qs[offset + 0] | (bq1_0->qs[offset + 1] << 8) | + (bq1_0->qs[offset + 2] << 16) | (bq1_0->qs[offset + 3] << 24); + + // Unpack 32 bits into 8 packed int32s (each holding 4 signed bytes: -1 or +1) + int vi_bytes[8]; +#pragma unroll + for (int j = 0; j < 8; ++j) { + const int shift = j * 4; + const int bits4 = (v >> shift) & 0x0F; + const int b0 = (bits4 & 0x01) ? 1 : -1; + const int b1 = (bits4 & 0x02) ? 1 : -1; + const int b2 = (bits4 & 0x04) ? 1 : -1; + const int b3 = (bits4 & 0x08) ? 1 : -1; + vi_bytes[j] = (b0 & 0xFF) | ((b1 & 0xFF) << 8) | ((b2 & 0xFF) << 16) | ((b3 & 0xFF) << 24); + } + + // dp4a dot product against Q8_1 activations + int sumi = 0; +#pragma unroll + for (int j = 0; j < 8; ++j) { + const int u = get_int_b4(bq8_1_chunk->qs, j); + sumi = __dp4a(vi_bytes[j], u, sumi); + } + + // Apply Q1_0's single scale and this chunk's Q8_1 scale + const float d8 = __low2float(bq8_1_chunk->ds); + return d1 * d8 * sumi; +#else + (void)vbq; (void)bq8_1; (void)iqs; + return 0.0f; +#endif +} + static __device__ __forceinline__ float vec_dot_q4_0_q8_1(const void* __restrict__ vbq, const block_q8_1* __restrict__ bq8_1, const int& iqs) { const block_q4_0* bq4_0 = (const block_q4_0*)vbq; @@ -1033,6 +1084,83 @@ static __device__ __forceinline__ void load_tiles_q8_0( } } +// Q1_0 MMQ support — unpack 1-bit to ±1 signed bytes, reuse Q8_0 dot product +// With QI1_0_MMQ=32, blocks_per_warp=1, each iteration = 1 Q1_0 block = 128 elements = 4 Q8_1 blocks +#define QI1_0_MMQ 32 + +template +static __device__ __forceinline__ void allocate_tiles_q1_0(int** x_ql, half2** x_dm, int** x_qh, int** x_sc) { + // Same layout as Q8_0: 32 int32s per row for qs, 4 floats per row for scales + __shared__ int tile_x_qs[mmq_y * (WARP_SIZE_GGUF) + mmq_y]; + __shared__ float tile_x_d[mmq_y * (WARP_SIZE_GGUF / QI8_0) + mmq_y / QI8_0]; + + *x_ql = tile_x_qs; + *x_dm = (half2*)tile_x_d; +} + +template +static __device__ __forceinline__ void load_tiles_q1_0( + const void* __restrict__ vx, + int* __restrict__ x_ql, + half2* __restrict__ x_dm, + int* __restrict__ x_qh, + int* __restrict__ x_sc, + const int& i_offset, + const int& i_max, + const int& k, + const int& blocks_per_row) { + + const block_q1_0* bx0 = (const block_q1_0*)vx; + float* x_dmf = (float*)x_dm; + + // k = threadIdx.x (0-31). Each thread unpacks 4 bits into 4 signed bytes. + // Elements k*4 through k*4+3 from the Q1_0 block. + const int base_elem = k * 4; // 0, 4, 8, ..., 124 + const int byte_idx = base_elem / 8; // 0..15 + const int bit_ofs = base_elem % 8; // always 0 or 4 + +#pragma unroll + for (int i0 = 0; i0 < mmq_y; i0 += nwarps) { + int i = i0 + i_offset; + + if (need_check) { + i = min(i, i_max); + } + + // With QI1_0_MMQ=32, blocks_per_warp=1, kbx=k/32=0 always + const block_q1_0* bxi = bx0 + i * blocks_per_row; + + const uint8_t byte_val = bxi->qs[byte_idx]; + const int bits4 = (byte_val >> bit_ofs) & 0x0F; + + // Unpack 4 bits to 4 signed bytes (-1 or +1) + const int b0 = (bits4 & 0x01) ? 1 : -1; + const int b1 = (bits4 & 0x02) ? 1 : -1; + const int b2 = (bits4 & 0x04) ? 1 : -1; + const int b3 = (bits4 & 0x08) ? 1 : -1; + + x_ql[i * (WARP_SIZE_GGUF + 1) + k] = + (b0 & 0xFF) | ((b1 & 0xFF) << 8) | ((b2 & 0xFF) << 16) | ((b3 & 0xFF) << 24); + } + + // Load scales — all 4 "virtual Q8_0 blocks" get the same Q1_0 scale d + const int blocks_per_tile_x_row = WARP_SIZE_GGUF / QI8_0; // = 4 + const int kbxd = k % blocks_per_tile_x_row; + +#pragma unroll + for (int i0 = 0; i0 < mmq_y; i0 += nwarps * QI8_0) { + int i = i0 + i_offset * QI8_0 + k / blocks_per_tile_x_row; + + if (need_check) { + i = min(i, i_max); + } + + // Same Q1_0 block for all 4 virtual positions (no kbxd offset) + const block_q1_0* bxi = bx0 + i * blocks_per_row; + x_dmf[i * (WARP_SIZE_GGUF / QI8_0) + i / QI8_0 + kbxd] = __half2float(bxi->d); + } +} + static __device__ __forceinline__ float vec_dot_q8_0_q8_1_mul_mat( const int* __restrict__ x_ql, const half2* __restrict__ x_dm,