diff --git a/native/kernels/dequant.cu b/native/kernels/dequant.cu
index a52a90fc..0ac4d773 100644
--- a/native/kernels/dequant.cu
+++ b/native/kernels/dequant.cu
@@ -107,6 +107,130 @@ extern "C" __global__ void __launch_bounds__(256) dequant_q5_0_f16(
     }
 }
 
+// ── Q4_1: 20 bytes per 32 values ────────────────────────────────────
+// struct block_q4_1 { half d; half m; uint8_t qs[16]; };
+// value = d * nibble + m  (unsigned + min, vs Q4_0's signed-after-bias-of-8)
+#define Q4_1_BLOCK_SIZE 32
+#define Q4_1_BLOCK_BYTES 20
+
+extern "C" __global__ void __launch_bounds__(256) dequant_q4_1_f16(
+    const uint8_t* __restrict__ src,
+    half* __restrict__ dst,
+    const int total_blocks)
+{
+    int lane = threadIdx.x % Q4_1_BLOCK_SIZE;
+    int warp_in_block = threadIdx.x / Q4_1_BLOCK_SIZE;
+    int warps_per_grid = (gridDim.x * blockDim.x) / Q4_1_BLOCK_SIZE;
+    int start_block = blockIdx.x * (blockDim.x / Q4_1_BLOCK_SIZE) + warp_in_block;
+
+    for (int block_idx = start_block; block_idx < total_blocks; block_idx += warps_per_grid)
+    {
+        const uint8_t* block = src + (size_t)block_idx * Q4_1_BLOCK_BYTES;
+        float d = __half2float(*reinterpret_cast<const half*>(block));
+        float m = __half2float(*reinterpret_cast<const half*>(block + 2));
+        const uint8_t* qs = block + 4;
+
+        int byte_idx = lane / 2;
+        uint8_t packed = qs[byte_idx];
+        int val = (lane & 1) ? (int)(packed >> 4) : (int)(packed & 0x0F);
+
+        dst[(size_t)block_idx * Q4_1_BLOCK_SIZE + lane] = __float2half(d * (float)val + m);
+    }
+}
+
+// ── Q5_1: 24 bytes per 32 values ────────────────────────────────────
+// struct block_q5_1 { half d; half m; uint32_t qh; uint8_t qs[16]; };
+// value = d * ((qh_bit << 4) | nibble) + m  (5-bit unsigned + min)
+#define Q5_1_BLOCK_SIZE 32
+#define Q5_1_BLOCK_BYTES 24
+
+extern "C" __global__ void __launch_bounds__(256) dequant_q5_1_f16(
+    const uint8_t* __restrict__ src,
+    half* __restrict__ dst,
+    const int total_blocks)
+{
+    int lane = threadIdx.x % Q5_1_BLOCK_SIZE;
+    int warp_in_block = threadIdx.x / Q5_1_BLOCK_SIZE;
+    int warps_per_grid = (gridDim.x * blockDim.x) / Q5_1_BLOCK_SIZE;
+    int start_block = blockIdx.x * (blockDim.x / Q5_1_BLOCK_SIZE) + warp_in_block;
+
+    for (int block_idx = start_block; block_idx < total_blocks; block_idx += warps_per_grid)
+    {
+        const uint8_t* block = src + (size_t)block_idx * Q5_1_BLOCK_BYTES;
+        float d = __half2float(*reinterpret_cast<const half*>(block));
+        float m = __half2float(*reinterpret_cast<const half*>(block + 2));
+        // Read qh (4 bytes, may be unaligned).
+        unsigned int qh = (unsigned int)block[4] | ((unsigned int)block[5] << 8) |
+                          ((unsigned int)block[6] << 16) | ((unsigned int)block[7] << 24);
+        const uint8_t* qs = block + 8;
+
+        int j = lane < 16 ? lane : lane - 16;
+        uint8_t packed = qs[j];
+        int nibble = (lane < 16) ? (packed & 0x0F) : (packed >> 4);
+        int high_bit = (qh >> lane) & 1;
+        int val = nibble | (high_bit << 4);
+
+        dst[(size_t)block_idx * Q5_1_BLOCK_SIZE + lane] = __float2half(d * (float)val + m);
+    }
+}
+
+// ── Q3_K: 110 bytes per 256 values (super-block with 16 sub-blocks of 16) ──
+// struct block_q3_K { uint8_t hmask[32]; uint8_t qs[64]; uint8_t scales[12]; half d; };
+//   hmask: 1 high bit per element (32 × 8 = 256 bits)
+//   qs:    2 low bits per element (64 × 4 = 256)
+//   scales: 16 × 6-bit signed-after-bias-of-32 packed into 12 bytes
+//   d: FP16 super-block delta
+// Per-element value: d × (signed_scale[sub]) × (((hmask_bit << 2) | qs_bits) - 4)
+#define Q3_K_SUPER_BLOCK_SIZE 256
+#define Q3_K_BLOCK_BYTES 110
+
+extern "C" __global__ void __launch_bounds__(256) dequant_q3_k_f16(
+    const uint8_t* __restrict__ src,
+    half* __restrict__ dst,
+    const int total_superblocks)
+{
+    int t = threadIdx.x; // 0..255
+
+    for (int sb_idx = blockIdx.x; sb_idx < total_superblocks; sb_idx += gridDim.x)
+    {
+        const uint8_t* block = src + (size_t)sb_idx * Q3_K_BLOCK_BYTES;
+        const uint8_t* hmask = block;                // 32 bytes
+        const uint8_t* qs = block + 32;              // 64 bytes
+        const uint8_t* scales12 = block + 32 + 64;   // 12 bytes
+        float d = __half2float(*reinterpret_cast<const half*>(block + 32 + 64 + 12));
+
+        // Sub-block index (0..15) for this thread; 16 threads share a sub-block
+        // (= 16 elements per sub-block).
+        int sub = t / 16;
+
+        // Unpack 6-bit scale for this sub-block from the 12 packed bytes.
+        // Per llama.cpp ggml-quants.c dequantize_row_q3_K:
+        //   sub 0..7  low nibble = scales12[sub] low nibble
+        //   sub 8..15 low nibble = scales12[sub-8] high nibble (NOT sub-4 — that
+        //   collides with the high-2-bits packing in scales12[8..11]).
+        //   high 2 bits = scales12[8 + (sub % 4)] >> ((sub / 4) * 2)
+        // The byte/shift pair is TRANSPOSED relative to the obvious-looking
+        // 8 + sub/4 @ (sub%4)*2 — the wrong way round scrambles 12 of 16 scales.
+        int lowSrcByte = sub < 8 ? sub : sub - 8;
+        int lowNibble = sub < 8
+            ? (scales12[lowSrcByte] & 0x0F)
+            : ((scales12[lowSrcByte] >> 4) & 0x0F);
+        int hiBits = (scales12[8 + (sub & 3)] >> ((sub >> 2) * 2)) & 0x03;
+        int signedScale = (lowNibble | (hiBits << 4)) - 32;  // [-32, 31]
+
+        // Per-element 3-bit unpacking. The 2-bit quants are NOT stored four
+        // consecutive elements per byte: element t reads bit-pair (t/32)%4 of
+        // qs byte (t%32) + 32*(t/128), and hmask bit t/32 of byte t%32
+        // (llama.cpp's shift/m loop over 128-element halves).
+        int qBits = (qs[(t & 31) + 32 * (t >> 7)] >> (((t >> 5) & 3) * 2)) & 0x03;
+        int hBit = (hmask[t & 31] >> (t >> 5)) & 0x01;
+        int signed3 = ((hBit << 2) | qBits) - 4;  // [-4, 3]
+
+        dst[(size_t)sb_idx * Q3_K_SUPER_BLOCK_SIZE + t] =
+            __float2half(d * (float)signedScale * (float)signed3);
+    }
+}
+
 // ── Q4_K: 144 bytes per 256 values (super-block with 8 sub-blocks) ──
 // struct block_q4_K { half d; half dmin; uint8_t scales[12]; uint8_t qs[128]; };
 #define Q4_K_SUPER_BLOCK_SIZE 256
diff --git a/native/ptx/dequant.ptx b/native/ptx/dequant.ptx
index 2f926da3..310b1193 100644
--- a/native/ptx/dequant.ptx
+++ b/native/ptx/dequant.ptx
@@ -1,12 +1,12 @@
 //
 // Generated by NVIDIA NVVM Compiler
 //
-// Compiler Build ID: CL-34714021
-// Cuda compilation tools, release 12.6, V12.6.68
+// Compiler Build ID: CL-35583870
+// Cuda compilation tools, release 12.8, V12.8.93
 // Based on NVVM 7.0.1
 //
 
-.version 8.5
+.version 8.7
 .target sm_61
 .address_size 64
 
@@ -60,7 +60,7 @@ $L__BB0_2:
 	ld.global.nc.u8 	%rs3, [%rd8+2];
 	cvt.s16.s8 	%rs4, %rs3;
 	cvt.rn.f32.s16 	%f3, %rs4;
-	mul.ftz.f32 	%f2, %f1, %f3;
+	mul.f32 	%f2, %f1, %f3;
 	// begin inline asm
 	{  cvt.rn.f16.f32 %rs2, %f2;}
 
@@ -133,7 +133,7 @@ $L__BB1_2:
 	mul.wide.s32 	%rd9, %r21, 32;
 	or.b64  	%rd10, %rd9, %rd2;
 	cvt.rn.f32.s32 	%f3, %r20;
-	mul.ftz.f32 	%f2, %f1, %f3;
+	mul.f32 	%f2, %f1, %f3;
 	// begin inline asm
 	{  cvt.rn.f16.f32 %rs2, %f2;}
 
@@ -223,7 +223,7 @@ $L__BB2_2:
 	mul.wide.s32 	%rd10, %r34, 32;
 	or.b64  	%rd11, %rd10, %rd2;
 	cvt.rn.f32.s32 	%f3, %r33;
-	mul.ftz.f32 	%f2, %f1, %f3;
+	mul.f32 	%f2, %f1, %f3;
 	// begin inline asm
 	{  cvt.rn.f16.f32 %rs2, %f2;}
 
@@ -238,6 +238,289 @@ $L__BB2_2:
 $L__BB2_3:
 	ret;
 
+}
+	// .globl	dequant_q4_1_f16
+.visible .entry dequant_q4_1_f16(
+	.param .u64 dequant_q4_1_f16_param_0,
+	.param .u64 dequant_q4_1_f16_param_1,
+	.param .u32 dequant_q4_1_f16_param_2
+)
+.maxntid 256, 1, 1
+{
+	.reg .pred 	%p<4>;
+	.reg .b16 	%rs<5>;
+	.reg .f32 	%f<5>;
+	.reg .b32 	%r<21>;
+	.reg .b64 	%rd<14>;
+
+
+	ld.param.u64 	%rd4, [dequant_q4_1_f16_param_0];
+	ld.param.u64 	%rd5, [dequant_q4_1_f16_param_1];
+	ld.param.u32 	%r8, [dequant_q4_1_f16_param_2];
+	mov.u32 	%r1, %tid.x;
+	shr.u32 	%r9, %r1, 5;
+	mov.u32 	%r2, %ntid.x;
+	shr.u32 	%r10, %r2, 5;
+	mov.u32 	%r11, %ctaid.x;
+	mad.lo.s32 	%r20, %r10, %r11, %r9;
+	setp.ge.s32 	%p1, %r20, %r8;
+	@%p1 bra 	$L__BB3_3;
+
+	and.b32  	%r12, %r1, 31;
+	bfe.u32 	%r13, %r1, 1, 4;
+	cvt.u64.u32 	%rd1, %r13;
+	and.b32  	%r4, %r1, 1;
+	cvt.u64.u32 	%rd2, %r12;
+	mov.u32 	%r14, %nctaid.x;
+	mul.lo.s32 	%r15, %r14, %r2;
+	shr.u32 	%r5, %r15, 5;
+	cvta.to.global.u64 	%rd3, %rd4;
+	cvta.to.global.u64 	%rd11, %rd5;
+
+$L__BB3_2:
+	mul.wide.s32 	%rd6, %r20, 20;
+	add.s64 	%rd7, %rd3, %rd6;
+	ld.global.nc.u16 	%rs1, [%rd7];
+	// begin inline asm
+	{  cvt.f32.f16 %f1, %rs1;}
+
+	// end inline asm
+	ld.global.nc.u16 	%rs2, [%rd7+2];
+	// begin inline asm
+	{  cvt.f32.f16 %f2, %rs2;}
+
+	// end inline asm
+	add.s64 	%rd8, %rd7, %rd1;
+	ld.global.nc.u8 	%rs4, [%rd8+4];
+	cvt.u32.u16 	%r16, %rs4;
+	bfe.u32 	%r17, %r16, 4, 4;
+	and.b32  	%r18, %r16, 15;
+	setp.eq.s32 	%p2, %r4, 0;
+	selp.b32 	%r19, %r18, %r17, %p2;
+	mul.wide.s32 	%rd9, %r20, 32;
+	or.b64  	%rd10, %rd9, %rd2;
+	cvt.rn.f32.s32 	%f4, %r19;
+	fma.rn.f32 	%f3, %f1, %f4, %f2;
+	// begin inline asm
+	{  cvt.rn.f16.f32 %rs3, %f3;}
+
+	// end inline asm
+	shl.b64 	%rd12, %rd10, 1;
+	add.s64 	%rd13, %rd11, %rd12;
+	st.global.u16 	[%rd13], %rs3;
+	add.s32 	%r20, %r20, %r5;
+	setp.lt.s32 	%p3, %r20, %r8;
+	@%p3 bra 	$L__BB3_2;
+
+$L__BB3_3:
+	ret;
+
+}
+	// .globl	dequant_q5_1_f16
+.visible .entry dequant_q5_1_f16(
+	.param .u64 dequant_q5_1_f16_param_0,
+	.param .u64 dequant_q5_1_f16_param_1,
+	.param .u32 dequant_q5_1_f16_param_2
+)
+.maxntid 256, 1, 1
+{
+	.reg .pred 	%p<5>;
+	.reg .b16 	%rs<9>;
+	.reg .f32 	%f<5>;
+	.reg .b32 	%r<34>;
+	.reg .b64 	%rd<14>;
+
+
+	ld.param.u64 	%rd5, [dequant_q5_1_f16_param_0];
+	ld.param.u64 	%rd6, [dequant_q5_1_f16_param_1];
+	ld.param.u32 	%r7, [dequant_q5_1_f16_param_2];
+	mov.u32 	%r8, %tid.x;
+	and.b32  	%r1, %r8, 31;
+	shr.u32 	%r9, %r8, 5;
+	mov.u32 	%r2, %ntid.x;
+	shr.u32 	%r10, %r2, 5;
+	mov.u32 	%r11, %ctaid.x;
+	mad.lo.s32 	%r33, %r10, %r11, %r9;
+	setp.ge.s32 	%p1, %r33, %r7;
+	@%p1 bra 	$L__BB4_3;
+
+	setp.lt.u32 	%p2, %r1, 16;
+	add.s32 	%r12, %r1, -16;
+	selp.b32 	%r13, %r1, %r12, %p2;
+	add.s32 	%r14, %r13, 8;
+	cvt.s64.s32 	%rd1, %r14;
+	cvt.u64.u32 	%rd2, %r1;
+	mov.u32 	%r15, %nctaid.x;
+	mul.lo.s32 	%r16, %r15, %r2;
+	shr.u32 	%r4, %r16, 5;
+	cvta.to.global.u64 	%rd3, %rd5;
+	cvta.to.global.u64 	%rd4, %rd6;
+
+$L__BB4_2:
+	mul.wide.s32 	%rd7, %r33, 24;
+	add.s64 	%rd8, %rd3, %rd7;
+	ld.global.nc.u16 	%rs1, [%rd8];
+	// begin inline asm
+	{  cvt.f32.f16 %f1, %rs1;}
+
+	// end inline asm
+	ld.global.nc.u16 	%rs2, [%rd8+2];
+	// begin inline asm
+	{  cvt.f32.f16 %f2, %rs2;}
+
+	// end inline asm
+	ld.global.nc.u8 	%rs4, [%rd8+4];
+	cvt.u32.u16 	%r17, %rs4;
+	and.b32  	%r18, %r17, 255;
+	ld.global.nc.u8 	%rs5, [%rd8+5];
+	cvt.u32.u16 	%r19, %rs5;
+	prmt.b32 	%r20, %r19, %r18, 30212;
+	ld.global.nc.u8 	%rs6, [%rd8+6];
+	cvt.u32.u16 	%r21, %rs6;
+	prmt.b32 	%r22, %r21, %r20, 28756;
+	ld.global.nc.u8 	%rs7, [%rd8+7];
+	cvt.u32.u16 	%r23, %rs7;
+	prmt.b32 	%r24, %r23, %r22, 1620;
+	add.s64 	%rd9, %rd8, %rd1;
+	ld.global.nc.u8 	%rs8, [%rd9];
+	cvt.u32.u16 	%r25, %rs8;
+	and.b32  	%r26, %r25, 15;
+	bfe.u32 	%r27, %r25, 4, 4;
+	selp.b32 	%r28, %r26, %r27, %p2;
+	shr.u32 	%r29, %r24, %r1;
+	shl.b32 	%r30, %r29, 4;
+	and.b32  	%r31, %r30, 16;
+	or.b32  	%r32, %r31, %r28;
+	mul.wide.s32 	%rd10, %r33, 32;
+	or.b64  	%rd11, %rd10, %rd2;
+	cvt.rn.f32.s32 	%f4, %r32;
+	fma.rn.f32 	%f3, %f1, %f4, %f2;
+	// begin inline asm
+	{  cvt.rn.f16.f32 %rs3, %f3;}
+
+	// end inline asm
+	shl.b64 	%rd12, %rd11, 1;
+	add.s64 	%rd13, %rd4, %rd12;
+	st.global.u16 	[%rd13], %rs3;
+	add.s32 	%r33, %r33, %r4;
+	setp.lt.s32 	%p4, %r33, %r7;
+	@%p4 bra 	$L__BB4_2;
+
+$L__BB4_3:
+	ret;
+
+}
+	// .globl	dequant_q3_k_f16
+.visible .entry dequant_q3_k_f16(
+	.param .u64 dequant_q3_k_f16_param_0,
+	.param .u64 dequant_q3_k_f16_param_1,
+	.param .u32 dequant_q3_k_f16_param_2
+)
+.maxntid 256, 1, 1
+{
+	.reg .pred 	%p<5>;
+	.reg .b16 	%rs<7>;
+	.reg .f32 	%f<6>;
+	.reg .b32 	%r<45>;
+	.reg .b64 	%rd<25>;
+
+
+	ld.param.u64 	%rd8, [dequant_q3_k_f16_param_0];
+	ld.param.u64 	%rd9, [dequant_q3_k_f16_param_1];
+	ld.param.u32 	%r9, [dequant_q3_k_f16_param_2];
+	mov.u32 	%r1, %tid.x;
+	mov.u32 	%r44, %ctaid.x;
+	setp.ge.s32 	%p1, %r44, %r9;
+	@%p1 bra 	$L__BB5_3;
+
+	cvta.to.global.u64 	%rd1, %rd8;
+	shr.s32 	%r10, %r1, 31;
+	shr.u32 	%r11, %r10, 28;
+	add.s32 	%r12, %r1, %r11;
+	shr.s32 	%r13, %r12, 4;
+	add.s32 	%r14, %r13, -4;
+	setp.lt.s32 	%p2, %r1, 128;
+	selp.b32 	%r15, %r13, %r14, %p2;
+	cvt.s64.s32 	%rd2, %r15;
+	shr.s32 	%r16, %r12, 6;
+	add.s32 	%r17, %r16, 8;
+	cvt.s64.s32 	%rd3, %r17;
+	shl.b32 	%r18, %r13, 1;
+	and.b32  	%r3, %r18, 6;
+	shr.s32 	%r19, %r1, 2;
+	add.s32 	%r20, %r19, 32;
+	cvt.s64.s32 	%rd4, %r20;
+	shl.b32 	%r21, %r1, 1;
+	and.b32  	%r4, %r21, 6;
+	shr.s32 	%r22, %r1, 3;
+	cvt.s64.s32 	%rd5, %r22;
+	and.b32  	%r5, %r1, 7;
+	cvt.s64.s32 	%rd6, %r1;
+	mov.u32 	%r6, %nctaid.x;
+	cvta.to.global.u64 	%rd7, %rd9;
+
+$L__BB5_2:
+	mul.wide.s32 	%rd10, %r44, 110;
+	add.s64 	%rd11, %rd10, 96;
+	add.s64 	%rd12, %rd1, %rd10;
+	ld.global.nc.u16 	%rs1, [%rd12+108];
+	// begin inline asm
+	{  cvt.f32.f16 %f1, %rs1;}
+
+	// end inline asm
+	add.s64 	%rd13, %rd11, %rd2;
+	add.s64 	%rd14, %rd1, %rd13;
+	ld.global.nc.u8 	%rs3, [%rd14];
+	cvt.u32.u16 	%r23, %rs3;
+	and.b32  	%r24, %r23, 15;
+	bfe.u32 	%r25, %r23, 4, 4;
+	selp.b32 	%r26, %r24, %r25, %p2;
+	add.s64 	%rd15, %rd11, %rd3;
+	add.s64 	%rd16, %rd1, %rd15;
+	ld.global.nc.u8 	%rs4, [%rd16];
+	cvt.u32.u16 	%r27, %rs4;
+	and.b32  	%r28, %r27, 255;
+	shr.u32 	%r29, %r28, %r3;
+	shl.b32 	%r30, %r29, 4;
+	and.b32  	%r31, %r30, 48;
+	or.b32  	%r32, %r31, %r26;
+	add.s32 	%r33, %r32, -32;
+	add.s64 	%rd17, %rd10, %rd4;
+	add.s64 	%rd18, %rd1, %rd17;
+	ld.global.nc.u8 	%rs5, [%rd18];
+	cvt.u32.u16 	%r34, %rs5;
+	and.b32  	%r35, %r34, 255;
+	shr.u32 	%r36, %r35, %r4;
+	and.b32  	%r37, %r36, 3;
+	add.s64 	%rd19, %rd10, %rd5;
+	add.s64 	%rd20, %rd1, %rd19;
+	ld.global.nc.u8 	%rs6, [%rd20];
+	cvt.u32.u16 	%r38, %rs6;
+	and.b32  	%r39, %r38, 255;
+	shr.u32 	%r40, %r39, %r5;
+	and.b32  	%r41, %r40, 1;
+	bfi.b32 	%r42, %r41, %r37, 2, 1;
+	add.s32 	%r43, %r42, -4;
+	mul.wide.s32 	%rd21, %r44, 256;
+	add.s64 	%rd22, %rd21, %rd6;
+	cvt.rn.f32.s32 	%f3, %r33;
+	mul.f32 	%f4, %f1, %f3;
+	cvt.rn.f32.s32 	%f5, %r43;
+	mul.f32 	%f2, %f4, %f5;
+	// begin inline asm
+	{  cvt.rn.f16.f32 %rs2, %f2;}
+
+	// end inline asm
+	shl.b64 	%rd23, %rd22, 1;
+	add.s64 	%rd24, %rd7, %rd23;
+	st.global.u16 	[%rd24], %rs2;
+	add.s32 	%r44, %r44, %r6;
+	setp.lt.s32 	%p4, %r44, %r9;
+	@%p4 bra 	$L__BB5_2;
+
+$L__BB5_3:
+	ret;
+
 }
 	// .globl	dequant_q4_k_f16
 .visible .entry dequant_q4_k_f16(
@@ -259,7 +542,7 @@ $L__BB2_3:
 	ld.param.u32 	%r6, [dequant_q4_k_f16_param_2];
 	mov.u32 	%r30, %ctaid.x;
 	setp.ge.s32 	%p1, %r30, %r6;
-	@%p1 bra 	$L__BB3_6;
+	@%p1 bra 	$L__BB6_6;
 
 	mov.u32 	%r7, %tid.x;
 	shr.s32 	%r8, %r7, 31;
@@ -285,7 +568,7 @@ $L__BB2_3:
 	add.s32 	%r3, %r20, -4;
 	cvta.to.global.u64 	%rd26, %rd4;
 
-$L__BB3_2:
+$L__BB6_2:
 	cvt.u32.u64 	%r23, %rd1;
 	setp.lt.s32 	%p3, %r23, 4;
 	mul.wide.s32 	%rd6, %r30, 144;
@@ -301,10 +584,10 @@ $L__BB3_2:
 	{  cvt.f32.f16 %f4, %rs8;}
 
 	// end inline asm
-	@%p3 bra 	$L__BB3_4;
-	bra.uni 	$L__BB3_3;
+	@%p3 bra 	$L__BB6_4;
+	bra.uni 	$L__BB6_3;
 
-$L__BB3_4:
+$L__BB6_4:
 	cvt.s64.s32 	%rd16, %r3;
 	add.s64 	%rd17, %rd6, %rd16;
 	add.s64 	%rd18, %rd7, %rd17;
@@ -312,9 +595,9 @@ $L__BB3_4:
 	and.b16  	%rs23, %rs19, 63;
 	ld.global.nc.u8 	%rs20, [%rd18+12];
 	and.b16  	%rs24, %rs20, 63;
-	bra.uni 	$L__BB3_5;
+	bra.uni 	$L__BB6_5;
 
-$L__BB3_3:
+$L__BB6_3:
 	cvt.s64.s32 	%rd11, %r3;
 	add.s64 	%rd12, %rd6, %rd11;
 	add.s64 	%rd13, %rd7, %rd12;
@@ -331,7 +614,7 @@ $L__BB3_3:
 	and.b16  	%rs18, %rs17, 48;
 	or.b16  	%rs24, %rs18, %rs15;
 
-$L__BB3_5:
+$L__BB6_5:
 	add.s64 	%rd22, %rd8, %rd2;
 	ld.global.nc.u8 	%rs22, [%rd22];
 	cvt.u32.u16 	%r24, %rs22;
@@ -339,12 +622,12 @@ $L__BB3_5:
 	and.b32  	%r26, %r24, 15;
 	selp.b32 	%r27, %r25, %r26, %p2;
 	cvt.rn.f32.u16 	%f6, %rs23;
-	mul.ftz.f32 	%f7, %f3, %f6;
+	mul.f32 	%f7, %f3, %f6;
 	cvt.rn.f32.s32 	%f8, %r27;
-	mul.ftz.f32 	%f9, %f7, %f8;
+	mul.f32 	%f9, %f7, %f8;
 	cvt.rn.f32.u16 	%f10, %rs24;
-	mul.ftz.f32 	%f11, %f4, %f10;
-	sub.ftz.f32 	%f5, %f9, %f11;
+	mul.f32 	%f11, %f4, %f10;
+	sub.f32 	%f5, %f9, %f11;
 	mul.wide.s32 	%rd23, %r30, 256;
 	cvt.s64.s32 	%rd24, %r7;
 	add.s64 	%rd25, %rd23, %rd24;
@@ -358,9 +641,9 @@ $L__BB3_5:
 	mov.u32 	%r29, %nctaid.x;
 	add.s32 	%r30, %r30, %r29;
 	setp.lt.s32 	%p5, %r30, %r6;
-	@%p5 bra 	$L__BB3_2;
+	@%p5 bra 	$L__BB6_2;
 
-$L__BB3_6:
+$L__BB6_6:
 	ret;
 
 }
@@ -384,7 +667,7 @@ $L__BB3_6:
 	ld.param.u32 	%r7, [dequant_q5_k_f16_param_2];
 	mov.u32 	%r45, %ctaid.x;
 	setp.ge.s32 	%p1, %r45, %r7;
-	@%p1 bra 	$L__BB4_6;
+	@%p1 bra 	$L__BB7_6;
 
 	mov.u32 	%r8, %tid.x;
 	shr.s32 	%r9, %r8, 31;
@@ -419,7 +702,7 @@ $L__BB3_6:
 	add.s32 	%r4, %r12, -4;
 	cvta.to.global.u64 	%rd25, %rd5;
 
-$L__BB4_2:
+$L__BB7_2:
 	setp.lt.s32 	%p2, %r8, 128;
 	mul.wide.s32 	%rd8, %r45, 176;
 	cvta.to.global.u64 	%rd9, %rd4;
@@ -434,10 +717,10 @@ $L__BB4_2:
 	{  cvt.f32.f16 %f4, %rs8;}
 
 	// end inline asm
-	@%p2 bra 	$L__BB4_4;
-	bra.uni 	$L__BB4_3;
+	@%p2 bra 	$L__BB7_4;
+	bra.uni 	$L__BB7_3;
 
-$L__BB4_4:
+$L__BB7_4:
 	cvt.s64.s32 	%rd17, %r4;
 	add.s64 	%rd18, %rd8, %rd17;
 	add.s64 	%rd19, %rd9, %rd18;
@@ -445,9 +728,9 @@ $L__BB4_4:
 	and.b16  	%rs24, %rs19, 63;
 	ld.global.nc.u8 	%rs20, [%rd19+12];
 	and.b16  	%rs25, %rs20, 63;
-	bra.uni 	$L__BB4_5;
+	bra.uni 	$L__BB7_5;
 
-$L__BB4_3:
+$L__BB7_3:
 	cvt.s64.s32 	%rd12, %r4;
 	add.s64 	%rd13, %rd8, %rd12;
 	add.s64 	%rd14, %rd9, %rd13;
@@ -464,11 +747,11 @@ $L__BB4_3:
 	and.b16  	%rs18, %rs17, 48;
 	or.b16  	%rs25, %rs18, %rs15;
 
-$L__BB4_5:
+$L__BB7_5:
 	cvt.rn.f32.u16 	%f6, %rs24;
-	mul.ftz.f32 	%f7, %f3, %f6;
+	mul.f32 	%f7, %f3, %f6;
 	cvt.rn.f32.u16 	%f8, %rs25;
-	mul.ftz.f32 	%f9, %f4, %f8;
+	mul.f32 	%f9, %f4, %f8;
 	add.s64 	%rd20, %rd3, %rd1;
 	ld.global.nc.u8 	%rs22, [%rd20];
 	cvt.u32.u16 	%r33, %rs22;
@@ -488,8 +771,8 @@ $L__BB4_5:
 	cvt.s64.s32 	%rd23, %r8;
 	add.s64 	%rd24, %rd22, %rd23;
 	cvt.rn.f32.s32 	%f10, %r42;
-	mul.ftz.f32 	%f11, %f7, %f10;
-	sub.ftz.f32 	%f5, %f11, %f9;
+	mul.f32 	%f11, %f7, %f10;
+	sub.f32 	%f5, %f11, %f9;
 	// begin inline asm
 	{  cvt.rn.f16.f32 %rs21, %f5;}
 
@@ -500,9 +783,9 @@ $L__BB4_5:
 	mov.u32 	%r44, %nctaid.x;
 	add.s32 	%r45, %r45, %r44;
 	setp.lt.s32 	%p4, %r45, %r7;
-	@%p4 bra 	$L__BB4_2;
+	@%p4 bra 	$L__BB7_2;
 
-$L__BB4_6:
+$L__BB7_6:
 	ret;
 
 }
@@ -526,7 +809,7 @@ $L__BB4_6:
 	ld.param.u32 	%r7, [dequant_q6_k_f16_param_2];
 	mov.u32 	%r42, %ctaid.x;
 	setp.ge.s32 	%p1, %r42, %r7;
-	@%p1 bra 	$L__BB5_10;
+	@%p1 bra 	$L__BB8_10;
 
 	mov.u32 	%r8, %tid.x;
 	shr.s32 	%r9, %r8, 31;
@@ -554,7 +837,7 @@ $L__BB4_6:
 	cvta.to.global.u64 	%rd8, %rd6;
 	cvta.to.global.u64 	%rd50, %rd7;
 
-$L__BB5_2:
+$L__BB8_2:
 	cvt.s64.s32 	%rd2, %r42;
 	mul.wide.s32 	%rd3, %r42, 210;
 	add.s64 	%rd9, %rd8, %rd3;
@@ -574,13 +857,13 @@ $L__BB5_2:
 	add.s64 	%rd15, %rd11, %rd14;
 	add.s64 	%rd5, %rd8, %rd15;
 	setp.eq.s32 	%p2, %r3, 0;
-	@%p2 bra 	$L__BB5_7;
+	@%p2 bra 	$L__BB8_7;
 
 	setp.eq.s32 	%p3, %r3, 1;
-	@%p3 bra 	$L__BB5_6;
+	@%p3 bra 	$L__BB8_6;
 
 	setp.ne.s32 	%p4, %r3, 2;
-	@%p4 bra 	$L__BB5_8;
+	@%p4 bra 	$L__BB8_8;
 
 	ld.global.nc.u8 	%rs7, [%rd4];
 	and.b16  	%rs8, %rs7, 240;
@@ -594,9 +877,9 @@ $L__BB5_2:
 	ld.global.nc.u8 	%rs10, [%rd22];
 	and.b16  	%rs11, %rs10, 48;
 	or.b16  	%rs31, %rs11, %rs9;
-	bra.uni 	$L__BB5_9;
+	bra.uni 	$L__BB8_9;
 
-$L__BB5_6:
+$L__BB8_6:
 	ld.global.nc.u8 	%rs12, [%rd5];
 	and.b16  	%rs13, %rs12, 15;
 	shl.b32 	%r31, %r2, 5;
@@ -609,9 +892,9 @@ $L__BB5_6:
 	shl.b16 	%rs15, %rs14, 2;
 	and.b16  	%rs16, %rs15, 48;
 	or.b16  	%rs31, %rs16, %rs13;
-	bra.uni 	$L__BB5_9;
+	bra.uni 	$L__BB8_9;
 
-$L__BB5_7:
+$L__BB8_7:
 	ld.global.nc.u8 	%rs17, [%rd4];
 	and.b16  	%rs18, %rs17, 15;
 	shl.b32 	%r33, %r2, 5;
@@ -624,9 +907,9 @@ $L__BB5_7:
 	shl.b16 	%rs20, %rs19, 4;
 	and.b16  	%rs21, %rs20, 48;
 	or.b16  	%rs31, %rs21, %rs18;
-	bra.uni 	$L__BB5_9;
+	bra.uni 	$L__BB8_9;
 
-$L__BB5_8:
+$L__BB8_8:
 	ld.global.nc.u8 	%rs22, [%rd5];
 	and.b16  	%rs23, %rs22, 240;
 	shr.u16 	%rs24, %rs23, 4;
@@ -641,7 +924,7 @@ $L__BB5_8:
 	and.b16  	%rs27, %rs26, 48;
 	or.b16  	%rs31, %rs27, %rs24;
 
-$L__BB5_9:
+$L__BB8_9:
 	cvt.u32.u16 	%r37, %rs31;
 	add.s32 	%r38, %r37, -32;
 	add.s64 	%rd45, %rd3, %rd1;
@@ -649,12 +932,12 @@ $L__BB5_9:
 	ld.global.nc.u8 	%rs29, [%rd46];
 	cvt.s16.s8 	%rs30, %rs29;
 	cvt.rn.f32.s16 	%f4, %rs30;
-	mul.ftz.f32 	%f5, %f2, %f4;
+	mul.f32 	%f5, %f2, %f4;
 	cvt.s64.s32 	%rd47, %r8;
 	shl.b64 	%rd48, %rd2, 8;
 	add.s64 	%rd49, %rd48, %rd47;
 	cvt.rn.f32.s32 	%f6, %r38;
-	mul.ftz.f32 	%f3, %f5, %f6;
+	mul.f32 	%f3, %f5, %f6;
 	// begin inline asm
 	{  cvt.rn.f16.f32 %rs28, %f3;}
 
@@ -666,9 +949,9 @@ $L__BB5_9:
 	cvt.u32.u64 	%r41, %rd2;
 	add.s32 	%r42, %r41, %r40;
 	setp.lt.s32 	%p5, %r42, %r7;
-	@%p5 bra 	$L__BB5_2;
+	@%p5 bra 	$L__BB8_2;
 
-$L__BB5_10:
+$L__BB8_10:
 	ret;
 
 }
diff --git a/src/DotLLM.Core/Configuration/QuantizationType.cs b/src/DotLLM.Core/Configuration/QuantizationType.cs
index fb9e4d19..6c078535 100644
--- a/src/DotLLM.Core/Configuration/QuantizationType.cs
+++ b/src/DotLLM.Core/Configuration/QuantizationType.cs
@@ -26,6 +26,12 @@ public enum QuantizationType
     /// <summary>8-bit quantization, group size 32.</summary>
     Q8_0 = 8,
 
+    /// <summary>2-bit K-quant, super-block of 256.</summary>
+    Q2_K = 10,
+
+    /// <summary>3-bit K-quant, super-block of 256.</summary>
+    Q3_K = 11,
+
     /// <summary>4-bit K-quant, super-block of 256.</summary>
     Q4_K = 12,
 
diff --git a/src/DotLLM.Core/Configuration/QuantizationTypeExtensions.cs b/src/DotLLM.Core/Configuration/QuantizationTypeExtensions.cs
index 374a2638..71642da1 100644
--- a/src/DotLLM.Core/Configuration/QuantizationTypeExtensions.cs
+++ b/src/DotLLM.Core/Configuration/QuantizationTypeExtensions.cs
@@ -20,6 +20,8 @@ public static class QuantizationTypeExtensions
         QuantizationType.Q5_0 => elementCount / 32 * 22,
         QuantizationType.Q5_1 => elementCount / 32 * 24,
         QuantizationType.Q8_0 => elementCount / 32 * 34,
+        QuantizationType.Q2_K => elementCount / 256 * 84,
+        QuantizationType.Q3_K => elementCount / 256 * 110,
         QuantizationType.Q4_K => elementCount / 256 * 144,
         QuantizationType.Q5_K => elementCount / 256 * 176,
         QuantizationType.Q6_K => elementCount / 256 * 210,
diff --git a/src/DotLLM.Cpu/Kernels/Dequantize.cs b/src/DotLLM.Cpu/Kernels/Dequantize.cs
index 63bb26f3..2a6faeae 100644
--- a/src/DotLLM.Cpu/Kernels/Dequantize.cs
+++ b/src/DotLLM.Cpu/Kernels/Dequantize.cs
@@ -29,6 +29,12 @@ public static unsafe partial class Dequantize
     /// <summary>Number of elements per Q5_0 block.</summary>
     private const int Q5_0GroupSize = 32;
 
+    /// <summary>Q4_1 block size in bytes: 2 (Half d) + 2 (Half m) + 16 (qs) = 20.</summary>
+    private const int Q4_1BlockBytes = 20;
+
+    /// <summary>Q5_1 block size in bytes: 2 (Half d) + 2 (Half m) + 4 (qh) + 16 (qs) = 24.</summary>
+    private const int Q5_1BlockBytes = 24;
+
     /// <summary>
     /// Returns the byte size of one row of <paramref name="elementCount"/> elements in the given quantization format.
     /// Useful for computing row strides when iterating weight matrices.
@@ -38,8 +44,12 @@ public static unsafe partial class Dequantize
         QuantizationType.F32 => elementCount * 4,
         QuantizationType.F16 => elementCount * 2,
         QuantizationType.Q4_0 => elementCount / Q8_0GroupSize * Q4_0BlockBytes,
+        QuantizationType.Q4_1 => elementCount / Q8_0GroupSize * Q4_1BlockBytes,
         QuantizationType.Q8_0 => elementCount / Q8_0GroupSize * Q8_0BlockBytes,
         QuantizationType.Q5_0 => elementCount / Q5_0GroupSize * Q5_0BlockBytes,
+        QuantizationType.Q5_1 => elementCount / Q5_0GroupSize * Q5_1BlockBytes,
+        QuantizationType.Q2_K => elementCount / KQuantGroupSize * Q2_K_BlockBytes,
+        QuantizationType.Q3_K => elementCount / KQuantGroupSize * Q3_K_BlockBytes,
         QuantizationType.Q4_K => elementCount / KQuantGroupSize * Q4_K_BlockBytes,
         QuantizationType.Q5_K => elementCount / KQuantGroupSize * Q5_K_BlockBytes,
         QuantizationType.Q6_K => elementCount / KQuantGroupSize * Q6_K_BlockBytes,
@@ -76,6 +86,18 @@ public static void ToFloat32(nint src, long elementCount, QuantizationType quant
             case QuantizationType.Q5_0:
                 DequantizeQ5_0(src, elementCount, dest);
                 break;
+            case QuantizationType.Q4_1:
+                DequantizeQ4_1Scalar(src, elementCount, dest);
+                break;
+            case QuantizationType.Q5_1:
+                DequantizeQ5_1Scalar(src, elementCount, dest);
+                break;
+            case QuantizationType.Q2_K:
+                DequantizeQ2_K(src, elementCount, dest);
+                break;
+            case QuantizationType.Q3_K:
+                DequantizeQ3_K(src, elementCount, dest);
+                break;
             case QuantizationType.Q4_K:
                 DequantizeQ4_K(src, elementCount, dest);
                 break;
@@ -149,6 +171,74 @@ internal static void DequantizeQ8_0Scalar(nint src, long elementCount, Span<floa
         }
     }
 
+    // ──────────────────── Q4_1 ────────────────────
+    /// <summary>
+    /// Q4_1 scalar dequant. Block layout (20 bytes, 32 elements):
+    /// <c>d(Half@0), m(Half@2), qs[16]@4</c>. Formula: <c>value = d * nibble + m</c>.
+    /// </summary>
+    [SkipLocalsInit]
+    internal static void DequantizeQ4_1Scalar(nint src, long elementCount, Span<float> dest)
+    {
+        if (elementCount % Q8_0GroupSize != 0)
+            throw new ArgumentException(
+                $"Q4_1 element count must be a multiple of {Q8_0GroupSize}, got {elementCount}",
+                nameof(elementCount));
+        long blockCount = elementCount / Q8_0GroupSize;
+        byte* blockBase = (byte*)src;
+        int outIdx = 0;
+        for (long b = 0; b < blockCount; b++)
+        {
+            float d = (float)Unsafe.ReadUnaligned<Half>(blockBase);
+            float m = (float)Unsafe.ReadUnaligned<Half>(blockBase + 2);
+            byte* qs = blockBase + 4;
+            for (int j = 0; j < 16; j++)
+            {
+                int lo = qs[j] & 0xF;
+                int hi = (qs[j] >> 4) & 0xF;
+                dest[outIdx + j]      = d * lo + m;
+                dest[outIdx + j + 16] = d * hi + m;
+            }
+            outIdx += Q8_0GroupSize;
+            blockBase += Q4_1BlockBytes;
+        }
+    }
+
+    // ──────────────────── Q5_1 ────────────────────
+    /// <summary>
+    /// Q5_1 scalar dequant. Block layout (24 bytes, 32 elements):
+    /// <c>d(Half@0), m(Half@2), qh[4]@4, qs[16]@8</c>.
+    /// Formula: <c>value = d * ((qh_bit &lt;&lt; 4) | nibble) + m</c> (5-bit unsigned + min).
+    /// </summary>
+    [SkipLocalsInit]
+    internal static void DequantizeQ5_1Scalar(nint src, long elementCount, Span<float> dest)
+    {
+        if (elementCount % Q5_0GroupSize != 0)
+            throw new ArgumentException(
+                $"Q5_1 element count must be a multiple of {Q5_0GroupSize}, got {elementCount}",
+                nameof(elementCount));
+        long blockCount = elementCount / Q5_0GroupSize;
+        byte* blockBase = (byte*)src;
+        int outIdx = 0;
+        for (long b = 0; b < blockCount; b++)
+        {
+            float d = (float)Unsafe.ReadUnaligned<Half>(blockBase);
+            float m = (float)Unsafe.ReadUnaligned<Half>(blockBase + 2);
+            uint qh = Unsafe.ReadUnaligned<uint>(blockBase + 4);
+            byte* qs = blockBase + 8;
+            for (int j = 0; j < 16; j++)
+            {
+                int lo = qs[j] & 0xF;
+                int hi = (qs[j] >> 4) & 0xF;
+                int bit5Lo = (int)((qh >> j) & 1);
+                int bit5Hi = (int)((qh >> (j + 16)) & 1);
+                dest[outIdx + j]      = d * (lo | (bit5Lo << 4)) + m;
+                dest[outIdx + j + 16] = d * (hi | (bit5Hi << 4)) + m;
+            }
+            outIdx += Q5_0GroupSize;
+            blockBase += Q5_1BlockBytes;
+        }
+    }
+
     // ──────────────────── Q5_0 ────────────────────
 
     [SkipLocalsInit]
diff --git a/src/DotLLM.Cpu/Kernels/DequantizeKQuants.cs b/src/DotLLM.Cpu/Kernels/DequantizeKQuants.cs
index 1f211782..e4e2b97d 100644
--- a/src/DotLLM.Cpu/Kernels/DequantizeKQuants.cs
+++ b/src/DotLLM.Cpu/Kernels/DequantizeKQuants.cs
@@ -11,6 +11,12 @@ namespace DotLLM.Cpu.Kernels;
 /// </summary>
 public static unsafe partial class Dequantize
 {
+    /// <summary>Q2_K block size in bytes: 16(scales) + 64(qs) + 2(d) + 2(dmin) = 84.</summary>
+    internal const int Q2_K_BlockBytes = 84;
+
+    /// <summary>Q3_K block size in bytes: 32(hmask) + 64(qs) + 12(scales) + 2(d) = 110.</summary>
+    internal const int Q3_K_BlockBytes = 110;
+
     /// <summary>Q4_K block size in bytes: 2(d) + 2(dmin) + 12(scales) + 128(qs) = 144.</summary>
     internal const int Q4_K_BlockBytes = 144;
 
@@ -208,6 +214,155 @@ internal static void DequantizeQ6_KAvx2(nint src, long elementCount, Span<float>
         }
     }
 
+    // ──────────────────── Q2_K ────────────────────
+
+    /// <summary>
+    /// Dequantizes Q2_K-quantized data to float32. Block layout:
+    /// scales[16] (4-bit scale + 4-bit dmin coef per sub-block, packed) +
+    /// qs[64] (2-bit elements, 4 per byte) + d (half) + dmin (half) = 84 bytes per 256 elements.
+    /// Per-element decode: <c>value = d × scale × q2 − dmin × dmin_coef</c>.
+    /// </summary>
+    [SkipLocalsInit]
+    public static unsafe void DequantizeQ2_K(nint src, long elementCount, Span<float> dest)
+    {
+        if (elementCount % KQuantGroupSize != 0)
+            throw new ArgumentException(
+                $"Q2_K requires elementCount to be a multiple of {KQuantGroupSize}.", nameof(elementCount));
+        // Bounds the int destination index below: dest.Length is an int, so this also
+        // guarantees elementCount <= int.MaxValue when called directly (not via ToFloat32).
+        if (dest.Length < elementCount)
+            throw new ArgumentException(
+                $"Destination span too small: {dest.Length} < {elementCount}", nameof(dest));
+
+        long superBlocks = elementCount / KQuantGroupSize;
+        byte* basePtr = (byte*)src;
+
+        for (long sb = 0; sb < superBlocks; sb++)
+        {
+            byte* block = basePtr + sb * Q2_K_BlockBytes;
+            byte* scales = block;          // 16 bytes
+            byte* qs = block + 16;         // 64 bytes
+            float d = (float)Unsafe.ReadUnaligned<Half>(block + 80);
+            float dmin = (float)Unsafe.ReadUnaligned<Half>(block + 82);
+
+            int outOffset = (int)(sb * KQuantGroupSize);
+            for (int t = 0; t < KQuantGroupSize; t++)
+            {
+                int sub = t >> 4;          // t / 16
+                int byteIdx = t >> 2;      // t / 4
+                int bitOff = (t & 0x3) << 1; // (t % 4) * 2
+                int q2 = (qs[byteIdx] >> bitOff) & 0x3;
+                int scale = scales[sub] & 0xF;
+                int dmCoef = (scales[sub] >> 4) & 0xF;
+                dest[outOffset + t] = d * scale * q2 - dmin * dmCoef;
+            }
+        }
+    }
+
+    // ──────────────────── Q3_K ────────────────────
+
+    /// <summary>
+    /// Dispatches Q3_K dequantization. Block layout (per ggml-quants.h):
+    /// <c>hmask[32]</c> (1 high bit per element) + <c>qs[64]</c> (2 low bits
+    /// per element) + <c>scales[12]</c> (16 packed 6-bit signed-after-bias
+    /// scales) + <c>d[2]</c> (FP16 super-block delta). 110 bytes per 256
+    /// elements. Per-element value:
+    /// <c>d × (signedScale[sub]) × ((hbit&lt;&lt;2 | qbits) - 4)</c> where sub
+    /// = element_idx / 16.
+    /// </summary>
+    [MethodImpl(MethodImplOptions.AggressiveInlining)]
+    internal static void DequantizeQ3_K(nint src, long elementCount, Span<float> dest)
+    {
+        if (elementCount % KQuantGroupSize != 0)
+            throw new ArgumentException(
+                $"Q3_K element count must be a multiple of {KQuantGroupSize}, got {elementCount}",
+                nameof(elementCount));
+        // Bounds the int destination index below: dest.Length is an int, so this also
+        // guarantees elementCount <= int.MaxValue when called directly (not via ToFloat32).
+        if (dest.Length < elementCount)
+            throw new ArgumentException(
+                $"Destination span too small: {dest.Length} < {elementCount}", nameof(dest));
+        DequantizeQ3_KScalar(src, elementCount, dest);
+    }
+
+    /// <summary>
+    /// Scalar Q3_K dequantization. Reference port of llama.cpp's
+    /// <c>dequantize_row_q3_K</c>. AVX2 acceleration is a future optimization;
+    /// at V2-Lite scale Q3_K is only used for token_embd.weight + output.weight
+    /// (not the per-call hot-path), so the scalar path is acceptable for now.
+    /// </summary>
+    internal static void DequantizeQ3_KScalar(nint src, long elementCount, Span<float> dest)
+    {
+        long numBlocks = elementCount / KQuantGroupSize;
+        byte* blockBase = (byte*)src;
+        // int base index, mirroring DequantizeQ2_K's outOffset — callers have already
+        // validated elementCount <= dest.Length, so this cannot overflow.
+        int destOffset = 0;
+        Span<byte> scales = stackalloc byte[16];
+
+        for (long b = 0; b < numBlocks; b++)
+        {
+            byte* hmask = blockBase;                 // [32 bytes]
+            byte* qs = blockBase + 32;               // [64 bytes]
+            byte* scales12 = blockBase + 32 + 64;    // [12 bytes]
+            ushort dHalf = *(ushort*)(blockBase + 32 + 64 + 12);
+            float d = (float)BitConverter.UInt16BitsToHalf(dHalf);
+
+            // Unpack 12 bytes → 16 unsigned 6-bit scales (then biased by -32).
+            // Per llama.cpp ggml-quants.c dequantize_row_q3_K (the `aux` shuffle):
+            //   scales[ 0+b] = (s[  b] & 0xF) | (((s[8+b] >> 0) & 3) << 4)
+            //   scales[ 4+b] = (s[4+b] & 0xF) | (((s[8+b] >> 2) & 3) << 4)
+            //   scales[ 8+b] = (s[  b] >> 4)  | (((s[8+b] >> 4) & 3) << 4)
+            //   scales[12+b] = (s[4+b] >> 4)  | (((s[8+b] >> 6) & 3) << 4)   for b in 0..3
+            // i.e. the low nibble comes from bytes 0..7 (low nibble for sub 0..7,
+            // high nibble for sub 8..15) and the high 2 bits come from
+            // byte 8 + (sub % 4) at shift 2 * (sub / 4). The byte/shift pair is
+            // TRANSPOSED relative to the obvious-looking 8 + sub/4 @ (sub%4)*2 —
+            // getting it the wrong way round scrambles 12 of the 16 sub-blocks.
+            for (int sub = 0; sub < 16; sub++)
+            {
+                int lowSrcByte = sub < 8 ? sub : sub - 8;  // sub 8..15 → bytes 0..7 high nibble
+                int lowNibble = sub < 8 ? scales12[lowSrcByte] & 0x0F : (scales12[lowSrcByte] >> 4) & 0x0F;
+                int hiByte = 8 + (sub % 4);
+                int hiShift = (sub / 4) * 2;
+                int hiBits = (scales12[hiByte] >> hiShift) & 0x03;
+                scales[sub] = (byte)(lowNibble | (hiBits << 4));
+            }
+
+            // 16 sub-blocks × 16 elements = 256 elements per super-block.
+            //
+            // Element ordering (llama.cpp dequantize_row_q3_K): the 2-bit quants
+            // are NOT stored 4-consecutive-elements-per-byte. Each 128-element
+            // half of the super-block uses 32 qs bytes, and each byte supplies
+            // FOUR elements 32 apart — element e reads bit-pair (e/32)%4 of
+            // byte (e%32) + 32*(e/128). The hmask is likewise transposed: element
+            // e reads bit e/32 of byte e%32. Reading them as e/4 @ (e%4)*2 and
+            // e/8 @ e%8 permutes every element of every super-block into the
+            // wrong sub-block scale — decoded weights correlate ~0.01 with the
+            // true values.
+            for (int sub = 0; sub < 16; sub++)
+            {
+                int signedScale = scales[sub] - 32;  // [-32, 31]
+                float scaleD = d * signedScale;
+                int eBase = sub * 16;
+                int qsSubBase = 32 * (sub >> 3) + 16 * (sub & 1);
+                int qShift = ((sub >> 1) & 3) * 2;
+                int hmSubBase = 16 * (sub & 1);
+                int hBitIdx = sub >> 1;
+                for (int l = 0; l < 16; l++)
+                {
+                    int qBits = (qs[qsSubBase + l] >> qShift) & 0x03;
+                    int hBit = (hmask[hmSubBase + l] >> hBitIdx) & 0x01;
+                    int signed3 = ((hBit << 2) | qBits) - 4;  // [-4, 3]
+                    dest[destOffset + eBase + l] = scaleD * signed3;
+                }
+            }
+
+            blockBase += Q3_K_BlockBytes;
+            destOffset += KQuantGroupSize;
+        }
+    }
+
     // ──────────────────── Q4_K ────────────────────
 
     /// <summary>Dispatches Q4_K dequantization.</summary>
diff --git a/src/DotLLM.Cuda/CudaKernels.cs b/src/DotLLM.Cuda/CudaKernels.cs
index ec4d059d..e353f2cc 100644
--- a/src/DotLLM.Cuda/CudaKernels.cs
+++ b/src/DotLLM.Cuda/CudaKernels.cs
@@ -77,7 +77,10 @@ public sealed unsafe class CudaKernels : IDisposable
     private readonly nint _quantizedGemvQ6_KFunc;
     private readonly nint _dequantQ8_0Func;
     private readonly nint _dequantQ4_0Func;
+    private readonly nint _dequantQ4_1Func;
     private readonly nint _dequantQ5_0Func;
+    private readonly nint _dequantQ5_1Func;
+    private readonly nint _dequantQ3_KFunc;
     private readonly nint _dequantQ4_KFunc;
     private readonly nint _dequantQ5_KFunc;
     private readonly nint _dequantQ6_KFunc;
@@ -150,7 +153,11 @@ public CudaKernels(string ptxDir)
         _quantizedGemvQ6_KFunc = _quantizedGemvModule.GetFunction("quantized_gemv_q6_k");
         _dequantQ8_0Func = _dequantModule.GetFunction("dequant_q8_0_f16");
         _dequantQ4_0Func = _dequantModule.GetFunction("dequant_q4_0_f16");
+        _dequantQ4_1Func = _dequantModule.TryGetFunction("dequant_q4_1_f16");
         _dequantQ5_0Func = _dequantModule.GetFunction("dequant_q5_0_f16");
+        _dequantQ5_1Func = _dequantModule.TryGetFunction("dequant_q5_1_f16");
+        // Q3_K is optional — older PTX builds (pre-Round 12) may not have it.
+        _dequantQ3_KFunc = _dequantModule.TryGetFunction("dequant_q3_k_f16");
         _dequantQ4_KFunc = _dequantModule.GetFunction("dequant_q4_k_f16");
         _dequantQ5_KFunc = _dequantModule.GetFunction("dequant_q5_k_f16");
         _dequantQ6_KFunc = _dequantModule.GetFunction("dequant_q6_k_f16");
@@ -650,6 +657,52 @@ public void LaunchDequantToF16(nint src, QuantizationType srcDtype,
                 return;
             }
 
+            case QuantizationType.Q4_1:
+            {
+                if (_dequantQ4_1Func == 0)
+                    throw new InvalidOperationException(
+                        "Q4_1 dequant kernel not in dequant.ptx — rebuild PTX from native/kernels/dequant.cu.");
+                int totalBlocks = totalElements / 32;
+                int tbArg = totalBlocks;
+                void** args = stackalloc void*[] {&srcArg, &dstArg, &tbArg};
+                uint gridDim = (uint)Math.Min((totalBlocks + 7) / 8, MaxDequantGridSize);
+                CudaDriverApi.cuLaunchKernel(_dequantQ4_1Func,
+                        gridDim, 1, 1, BlockSize, 1, 1,
+                        0, stream, (nint)args, 0).ThrowOnError();
+                return;
+            }
+
+            case QuantizationType.Q5_1:
+            {
+                if (_dequantQ5_1Func == 0)
+                    throw new InvalidOperationException(
+                        "Q5_1 dequant kernel not in dequant.ptx — rebuild PTX from native/kernels/dequant.cu.");
+                int totalBlocks = totalElements / 32;
+                int tbArg = totalBlocks;
+                void** args = stackalloc void*[] {&srcArg, &dstArg, &tbArg};
+                uint gridDim = (uint)Math.Min((totalBlocks + 7) / 8, MaxDequantGridSize);
+                CudaDriverApi.cuLaunchKernel(_dequantQ5_1Func,
+                        gridDim, 1, 1, BlockSize, 1, 1,
+                        0, stream, (nint)args, 0).ThrowOnError();
+                return;
+            }
+
+            case QuantizationType.Q3_K:
+            {
+                if (_dequantQ3_KFunc == 0)
+                    throw new InvalidOperationException(
+                        "Q3_K dequant kernel not present in dequant.ptx — rebuild PTX from " +
+                        "native/kernels/dequant.cu (Round 12+ adds Q3_K support).");
+                int totalSuperblocks = totalElements / 256;
+                int tsbArg = totalSuperblocks;
+                void** args = stackalloc void*[] {&srcArg, &dstArg, &tsbArg};
+                uint gridDim = (uint)Math.Min(totalSuperblocks, MaxDequantGridSize);
+                CudaDriverApi.cuLaunchKernel(_dequantQ3_KFunc,
+                        gridDim, 1, 1, BlockSize, 1, 1,
+                        0, stream, (nint)args, 0).ThrowOnError();
+                return;
+            }
+
             case QuantizationType.Q4_K:
             {
                 int totalSuperblocks = totalElements / 256;
diff --git a/src/DotLLM.Cuda/CudaModule.cs b/src/DotLLM.Cuda/CudaModule.cs
index 734ac424..d01ab5c7 100644
--- a/src/DotLLM.Cuda/CudaModule.cs
+++ b/src/DotLLM.Cuda/CudaModule.cs
@@ -12,6 +12,10 @@ public sealed class CudaModule : IDisposable
     private nint _module;
     private readonly Dictionary<string, nint> _functions = new();
 
+    // Misses are tracked separately from _functions: caching a miss as 0 in _functions
+    // would make GetFunction return 0 for that name instead of throwing.
+    private readonly HashSet<string> _missingFunctions = new();
+
     /// <summary>
     /// Loads a PTX module from a file path.
     /// </summary>
@@ -64,6 +68,35 @@ public nint GetFunction(string name)
         return func;
     }
 
+    /// <summary>
+    /// Tries to get a kernel function handle by name. Returns <c>0</c> if the symbol is not
+    /// present in the module (e.g. compiled against an older PTX that predates the kernel).
+    /// Caches both hits and misses for subsequent calls.
+    /// </summary>
+    /// <param name="name">The <c>extern "C"</c> kernel function name.</param>
+    /// <returns>The function handle, or <c>0</c> if the symbol was not found.</returns>
+    public nint TryGetFunction(string name)
+    {
+        if (_functions.TryGetValue(name, out nint func))
+            return func;
+        if (_missingFunctions.Contains(name))
+            return 0;
+
+        int result = CudaDriverApi.cuModuleGetFunction(out func, _module, name);
+
+        // Symbol absent from PTX (older build) — record it in the miss set, never in
+        // _functions, so a later GetFunction(name) still performs a real lookup and throws.
+        if (result == CudaResult.NotFound)
+        {
+            _missingFunctions.Add(name);
+            return 0;
+        }
+
+        result.ThrowOnError();
+        _functions[name] = func;
+        return func;
+    }
+
 
     /// <inheritdoc/>
     public void Dispose()
@@ -73,6 +106,7 @@ public void Dispose()
         {
             CudaDriverApi.cuModuleUnload(module);
             _functions.Clear();
+            _missingFunctions.Clear();
         }
     }
 }
diff --git a/src/DotLLM.Cuda/Interop/CudaErrorHelper.cs b/src/DotLLM.Cuda/Interop/CudaErrorHelper.cs
index e128a623..11c0d0de 100644
--- a/src/DotLLM.Cuda/Interop/CudaErrorHelper.cs
+++ b/src/DotLLM.Cuda/Interop/CudaErrorHelper.cs
@@ -2,6 +2,22 @@
 
 namespace DotLLM.Cuda.Interop;
 
+/// <summary>
+/// Named CUDA driver API result codes (<c>CUresult</c>) that call sites branch on,
+/// rather than comparing against bare magic numbers.
+/// </summary>
+internal static class CudaResult
+{
+    /// <summary><c>CUDA_SUCCESS</c>.</summary>
+    internal const int Success = 0;
+
+    /// <summary>
+    /// <c>CUDA_ERROR_NOT_FOUND</c> — a named symbol (kernel, global, texture) does not
+    /// exist in the loaded module.
+    /// </summary>
+    internal const int NotFound = 500;
+}
+
 /// <summary>
 /// Extension methods for checking CUDA and cuBLAS return codes.
 /// </summary>
diff --git a/tests/DotLLM.Tests.Unit/Cpu/Kernels/DequantizeKQuantTests.cs b/tests/DotLLM.Tests.Unit/Cpu/Kernels/DequantizeKQuantTests.cs
index 2651ca03..0e1ea4e3 100644
--- a/tests/DotLLM.Tests.Unit/Cpu/Kernels/DequantizeKQuantTests.cs
+++ b/tests/DotLLM.Tests.Unit/Cpu/Kernels/DequantizeKQuantTests.cs
@@ -202,6 +202,309 @@ public void Q6_K_MultipleBlocks(int blockCount)
         }
     }
 
+    // ──────────────────── Q3_K dequant ────────────────────
+
+    private const int Q3_K_BlockBytes = 110;
+
+    [Fact]
+    public void Q3_K_SingleBlock_HandCalculated()
+    {
+        // Block layout: hmask[32] + qs[64] + scales[12] + d[2] = 110 bytes.
+        nuint totalBytes = Q3_K_BlockBytes;
+        nint ptr = (nint)NativeMemory.AlignedAlloc(totalBytes, 64);
+        try
+        {
+            NativeMemory.Clear((void*)ptr, totalBytes);
+            byte* block = (byte*)ptr;
+
+            // d = 1.0
+            Unsafe.WriteUnaligned(block + 32 + 64 + 12, (Half)1.0f);
+
+            // scales12 (offset 32+64=96):
+            //   sub 0 → low nibble in scales12[0] (low 4 bits) + high 2 bits in scales12[8] bits 0-1
+            //   We want unsigned scale = 33 (= 32 + 1 → signed scale = +1).
+            //   33 = 0b100001 → low nibble 0b0001 (=1), high 2 bits 0b10 (=2).
+            block[96 + 0] = 0x01;        // scales12[0] = low nibble
+            block[96 + 8] = 0x02;        // scales12[8] bit 0-1 = high 2 bits of scale[0]
+
+            // qs[0] (offset 32): set element 0's 2 low bits to 0b11 (= 3)
+            block[32 + 0] = 0x03;
+
+            // hmask[0] (offset 0): set element 0's high bit to 1
+            block[0] = 0x01;
+
+            // Element 0: signed_3bit = ((1<<2) | 3) - 4 = 7 - 4 = 3
+            // Signed scale = 33 - 32 = 1
+            // d × scale × signed_3bit = 1.0 × 1 × 3 = 3.0
+            float[] dest = new float[KQuantGroupSize];
+            Dequantize.ToFloat32(ptr, KQuantGroupSize, QuantizationType.Q3_K, dest);
+
+            Assert.Equal(3.0f, dest[0], 0.01f);
+
+            // Element 1 (no qs/hmask bits set, scale[0] = 1):
+            //   signed_3bit = (0 << 2 | 0) - 4 = -4
+            //   value = 1.0 × 1 × -4 = -4
+            Assert.Equal(-4.0f, dest[1], 0.01f);
+
+            // Sub-block 1 (elements 16..31) has scale[1] = 0 - 32 = -32 → all values = 1 × -32 × -4 = 128
+            // (since qs/hmask are all zero, signed_3bit = -4 for every element).
+            Assert.Equal(128.0f, dest[16], 0.01f);
+        }
+        finally
+        {
+            NativeMemory.AlignedFree((void*)ptr);
+        }
+    }
+
+    /// <summary>
+    /// Discriminating Q3_K oracle test (#311). <see cref="Q3_K_SingleBlock_HandCalculated"/>
+    /// only touches element 0/1 of sub-block 0 and an all-zero sub-block 1 — a degenerate
+    /// case where the correct and the incorrect bit layouts coincide, which is exactly why
+    /// a total scramble of Q3_K could sit in this PR undetected.
+    ///
+    /// This test drives DENSE pseudorandom super-block bytes through
+    /// <see cref="Dequantize.ToFloat32"/> and compares against a LITERAL transcription of
+    /// llama.cpp's <c>ggml-quants.c dequantize_row_q3_K</c> — including the 32-bit
+    /// <c>aux</c>/<c>kmask</c> scale shuffle and the <c>shift</c>/<c>m</c> loop over
+    /// 128-element halves. The reference is written in llama.cpp's own control-flow shape,
+    /// structurally unlike the production kernel's closed-form indexing, so agreement is
+    /// evidence rather than a shared-mistake tautology.
+    ///
+    /// Discrimination proof: reverting either half of the fix makes this red —
+    /// the scale hi-bits byte/shift transposition (<c>8 + sub%4 @ (sub/4)*2</c> →
+    /// <c>8 + sub/4 @ (sub%4)*2</c>) or the element ordering
+    /// (<c>qs[(t%32)+32*(t/128)] @ ((t/32)%4)*2</c> → <c>qs[t/4] @ (t%4)*2</c>).
+    /// </summary>
+    [Fact]
+    public void Q3_K_DenseRandomBlocks_MatchLlamaCppReference()
+    {
+        const int blocks = 5;
+        const int elements = blocks * KQuantGroupSize;
+        nuint totalBytes = (nuint)(blocks * Q3_K_BlockBytes);
+        nint ptr = (nint)NativeMemory.AlignedAlloc(totalBytes, 64);
+        try
+        {
+            var rng = new Random(20260810);
+            byte* raw = (byte*)ptr;
+            for (int i = 0; i < (int)totalBytes; i++) raw[i] = (byte)rng.Next(256);
+            // Keep the fp16 super-block deltas finite and O(1) so the comparison is
+            // about bit layout, not about NaN/Inf plumbing.
+            for (int b = 0; b < blocks; b++)
+                Unsafe.WriteUnaligned(raw + b * Q3_K_BlockBytes + 108, (Half)(0.25f + 0.125f * b));
+
+            float[] actual = new float[elements];
+            Dequantize.ToFloat32(ptr, elements, QuantizationType.Q3_K, actual);
+
+            float[] expected = LlamaCppDequantizeRowQ3K(raw, blocks);
+
+            // Both sides compute the identical product in float — require exact equality.
+            for (int i = 0; i < elements; i++)
+            {
+                Assert.True(expected[i] == actual[i],
+                    $"Q3_K element {i} (block {i / KQuantGroupSize}, sub {(i % KQuantGroupSize) / 16}, "
+                    + $"lane {i % 16}): llama.cpp reference {expected[i]} != dotLLM {actual[i]}");
+            }
+        }
+        finally
+        {
+            NativeMemory.AlignedFree((void*)ptr);
+        }
+    }
+
+    /// <summary>
+    /// Literal transcription of llama.cpp <c>ggml-quants.c dequantize_row_q3_K</c>
+    /// (the authoritative GGUF Q3_K semantics), kept in its original control-flow shape
+    /// on purpose — see <see cref="Q3_K_DenseRandomBlocks_MatchLlamaCppReference"/>.
+    /// </summary>
+    private static float[] LlamaCppDequantizeRowQ3K(byte* src, int nb)
+    {
+        const uint kmask1 = 0x03030303u;
+        const uint kmask2 = 0x0f0f0f0fu;
+
+        var y = new float[nb * KQuantGroupSize];
+        int outIdx = 0;
+        uint* aux = stackalloc uint[4];
+        sbyte* scales = (sbyte*)aux;
+
+        for (int i = 0; i < nb; i++)
+        {
+            byte* block = src + i * Q3_K_BlockBytes;
+            byte* hm = block;               // hmask[32]
+            byte* q = block + 32;           // qs[64]
+            float dAll = (float)Unsafe.ReadUnaligned<Half>(block + 108);
+
+            for (int w = 0; w < 3; w++) aux[w] = Unsafe.ReadUnaligned<uint>(block + 96 + w * 4);
+            uint tmp = aux[2];
+            aux[2] = ((aux[0] >> 4) & kmask2) | (((tmp >> 4) & kmask1) << 4);
+            aux[3] = ((aux[1] >> 4) & kmask2) | (((tmp >> 6) & kmask1) << 4);
+            aux[0] = (aux[0] & kmask2) | (((tmp >> 0) & kmask1) << 4);
+            aux[1] = (aux[1] & kmask2) | (((tmp >> 2) & kmask1) << 4);
+
+            byte m = 1;
+            int qOff = 0;
+            int isIdx = 0;
+            for (int n = 0; n < KQuantGroupSize; n += 128)
+            {
+                int shift = 0;
+                for (int j = 0; j < 4; ++j)
+                {
+                    float dl = dAll * (scales[isIdx++] - 32);
+                    for (int l = 0; l < 16; ++l)
+                        y[outIdx++] = dl * (((q[qOff + l] >> shift) & 3) - (((hm[l] & m) != 0) ? 0 : 4));
+
+                    dl = dAll * (scales[isIdx++] - 32);
+                    for (int l = 0; l < 16; ++l)
+                        y[outIdx++] = dl * (((q[qOff + l + 16] >> shift) & 3) - (((hm[l + 16] & m) != 0) ? 0 : 4));
+
+                    shift += 2;
+                    m <<= 1;
+                }
+                qOff += 32;
+            }
+        }
+        return y;
+    }
+
+    [Fact]
+    public void Q3_K_RowByteSize_Matches()
+    {
+        // 256 elements = 1 super-block = 110 bytes.
+        Assert.Equal(110L, Dequantize.RowByteSize(256, QuantizationType.Q3_K));
+        // 1024 elements = 4 super-blocks = 440 bytes.
+        Assert.Equal(440L, Dequantize.RowByteSize(1024, QuantizationType.Q3_K));
+    }
+
+    [Fact]
+    public void Q3_K_NonAlignedCount_Throws()
+    {
+        float[] dest = new float[100];
+        Assert.Throws<ArgumentException>(() =>
+            Dequantize.ToFloat32(nint.Zero, 100, QuantizationType.Q3_K, dest));
+    }
+
+    // ──────────────────── Q2_K dequant ────────────────────
+
+    private const int Q2_K_BlockBytes = 84;
+
+    [Fact]
+    public void Q2_K_SingleBlock_HandCalculated()
+    {
+        // Block layout: scales[16] + qs[64] + d[2] + dmin[2] = 84 bytes.
+        nuint totalBytes = Q2_K_BlockBytes;
+        nint ptr = (nint)NativeMemory.AlignedAlloc(totalBytes, 64);
+        try
+        {
+            NativeMemory.Clear((void*)ptr, totalBytes);
+            byte* block = (byte*)ptr;
+
+            // d = 1.0, dmin = 0.5
+            Unsafe.WriteUnaligned(block + 80, (Half)1.0f);
+            Unsafe.WriteUnaligned(block + 82, (Half)0.5f);
+
+            // scales[0]: low nibble = scale (we want scale = 3), high nibble = dmin coef (we want 2).
+            // Packed as: (dmin_coef << 4) | scale = (2 << 4) | 3 = 0x23
+            block[0] = 0x23;
+
+            // qs[0] (offset 16): set element 0's 2 low bits to 0b10 (= 2).
+            // qs encoding: 4 elements per byte, low-to-high.
+            //   byte 0, bits 0-1 → element 0
+            //   byte 0, bits 2-3 → element 1
+            //   byte 0, bits 4-5 → element 2
+            //   byte 0, bits 6-7 → element 3
+            block[16 + 0] = 0x02;  // element 0 = 2, elements 1-3 = 0
+
+            // Element 0: q2 = 2, scale = 3, dmin_coef = 2
+            //   value = d * scale * q2 - dmin * dmin_coef
+            //         = 1.0 * 3 * 2 - 0.5 * 2
+            //         = 6 - 1 = 5
+            float[] dest = new float[KQuantGroupSize];
+            Dequantize.ToFloat32(ptr, KQuantGroupSize, QuantizationType.Q2_K, dest);
+
+            Assert.Equal(5.0f, dest[0], 0.01f);
+
+            // Element 1: q2 = 0, scale = 3, dmin_coef = 2
+            //   value = 1.0 * 3 * 0 - 0.5 * 2 = -1
+            Assert.Equal(-1.0f, dest[1], 0.01f);
+
+            // Sub-block 1 (elements 16..31): scale = 0, dmin_coef = 0 (all-zero scales[1..15])
+            //   value = 1.0 * 0 * 0 - 0.5 * 0 = 0
+            Assert.Equal(0.0f, dest[16], 0.01f);
+        }
+        finally
+        {
+            NativeMemory.AlignedFree((void*)ptr);
+        }
+    }
+
+    [Fact]
+    public void Q2_K_RowByteSize_Matches()
+    {
+        // 256 elements = 1 super-block = 84 bytes.
+        Assert.Equal(84L, Dequantize.RowByteSize(256, QuantizationType.Q2_K));
+        // 1024 elements = 4 super-blocks = 336 bytes.
+        Assert.Equal(336L, Dequantize.RowByteSize(1024, QuantizationType.Q2_K));
+    }
+
+    [Fact]
+    public void Q2_K_ComputeByteCount_MatchesRowByteSize()
+    {
+        // QuantizationTypeExtensions must agree with the CPU block size (84 bytes / 256).
+        Assert.Equal(84L, QuantizationType.Q2_K.ComputeByteCount(256));
+        Assert.Equal(336L, QuantizationType.Q2_K.ComputeByteCount(1024));
+        Assert.Equal(
+            Dequantize.RowByteSize(1024, QuantizationType.Q2_K),
+            QuantizationType.Q2_K.ComputeByteCount(1024));
+    }
+
+    [Fact]
+    public void Q2_K_NonAlignedCount_Throws()
+    {
+        float[] dest = new float[100];
+        Assert.Throws<ArgumentException>(() =>
+            Dequantize.ToFloat32(nint.Zero, 100, QuantizationType.Q2_K, dest));
+    }
+
+    [Fact]
+    public void Q2_K_TwoSuperBlocks_StrideCorrect()
+    {
+        // Two super-blocks of 256 elements each = 168 bytes total.
+        // SB0: d=1.0, dmin=0.0, scales[0]=0x03 (scale=3, dmin_coef=0), qs[0]=0x01 (element 0 q2=1)
+        // SB1: d=2.0, dmin=0.0, scales[0]=0x05 (scale=5, dmin_coef=0), qs[0]=0x03 (element 0 q2=3)
+        // Expect: dest[0]   = 1.0 * 3 * 1 - 0 = 3.0       (SB0, element 0)
+        //         dest[256] = 2.0 * 5 * 3 - 0 = 30.0      (SB1, element 0)
+        // Catches super-block stride bugs (e.g. sb*80 instead of sb*84).
+        nuint totalBytes = 2 * Q2_K_BlockBytes;  // 168
+        nint ptr = (nint)NativeMemory.AlignedAlloc(totalBytes, 64);
+        try
+        {
+            NativeMemory.Clear((void*)ptr, totalBytes);
+            byte* sb0 = (byte*)ptr;
+            byte* sb1 = (byte*)ptr + Q2_K_BlockBytes;
+
+            // SB0
+            Unsafe.WriteUnaligned(sb0 + 80, (Half)1.0f);
+            Unsafe.WriteUnaligned(sb0 + 82, (Half)0.0f);
+            sb0[0] = 0x03;
+            sb0[16] = 0x01;
+
+            // SB1
+            Unsafe.WriteUnaligned(sb1 + 80, (Half)2.0f);
+            Unsafe.WriteUnaligned(sb1 + 82, (Half)0.0f);
+            sb1[0] = 0x05;
+            sb1[16] = 0x03;
+
+            float[] dest = new float[2 * KQuantGroupSize];
+            Dequantize.ToFloat32(ptr, 2 * KQuantGroupSize, QuantizationType.Q2_K, dest);
+
+            Assert.Equal(3.0f,  dest[0],   0.01f);
+            Assert.Equal(30.0f, dest[256], 0.01f);
+        }
+        finally
+        {
+            NativeMemory.AlignedFree((void*)ptr);
+        }
+    }
+
     // ──────────────────── Q4_K dequant ────────────────────
 
     [Fact]
diff --git a/tests/DotLLM.Tests.Unit/Cpu/Kernels/DequantizeTests.cs b/tests/DotLLM.Tests.Unit/Cpu/Kernels/DequantizeTests.cs
index e80a3d74..d99f68d8 100644
--- a/tests/DotLLM.Tests.Unit/Cpu/Kernels/DequantizeTests.cs
+++ b/tests/DotLLM.Tests.Unit/Cpu/Kernels/DequantizeTests.cs
@@ -12,6 +12,8 @@ public sealed unsafe class DequantizeTests
     private const int Q8_0GroupSize = 32;
     private const int Q5_0BlockBytes = 22;
     private const int Q5_0GroupSize = 32;
+    private const int Q4_1BlockBytes = 20;
+    private const int Q5_1BlockBytes = 24;
 
     // ──────────────────── FP16 ────────────────────
 
@@ -375,8 +377,207 @@ public void Q8_0_NonAlignedCount_Throws()
             Dequantize.ToFloat32(nint.Zero, 33, QuantizationType.Q8_0, dest));
     }
 
+    // ──────────────────── Q4_1 ────────────────────
+
+    [Fact]
+    public void Q4_1_SingleBlock_HandCalculated()
+    {
+        // Layout (20 bytes / 32 elements): d(Half@0), m(Half@2), qs[16]@4.
+        // value = d * nibble + m, with the low nibble of qs[j] feeding element j and
+        // the high nibble feeding element j + 16 (llama.cpp dequantize_row_q4_1).
+        // d = 0.5, m = -3. qs[0] = 0x9A → lo = 0xA (10), hi = 0x9 (9).
+        //   dest[0]  = 0.5 * 10 - 3 = 2.0
+        //   dest[16] = 0.5 *  9 - 3 = 1.5
+        // qs[15] = 0x0F → lo = 15, hi = 0.
+        //   dest[15] = 0.5 * 15 - 3 = 4.5
+        //   dest[31] = 0.5 *  0 - 3 = -3.0
+        nint ptr = AllocQ4_1Block((Half)0.5f, (Half)(-3.0f), j => j switch
+        {
+            0 => (byte)0x9A,
+            15 => (byte)0x0F,
+            _ => (byte)0x00,
+        });
+        try
+        {
+            float[] dest = new float[Q8_0GroupSize];
+            Dequantize.ToFloat32(ptr, Q8_0GroupSize, QuantizationType.Q4_1, dest);
+
+            Assert.Equal(2.0f, dest[0], 1e-4f);
+            Assert.Equal(1.5f, dest[16], 1e-4f);
+            Assert.Equal(4.5f, dest[15], 1e-4f);
+            Assert.Equal(-3.0f, dest[31], 1e-4f);
+            // Every untouched nibble decodes to the block minimum.
+            for (int j = 1; j < 15; j++)
+            {
+                Assert.Equal(-3.0f, dest[j], 1e-4f);
+                Assert.Equal(-3.0f, dest[j + 16], 1e-4f);
+            }
+        }
+        finally
+        {
+            NativeMemory.AlignedFree((void*)ptr);
+        }
+    }
+
+    [Fact]
+    public void Q4_1_TwoBlocks_StrideCorrect()
+    {
+        // Catches a wrong block stride (e.g. 18 bytes — the Q4_0 size — instead of 20).
+        const int blockCount = 2;
+        nint ptr = (nint)NativeMemory.AlignedAlloc(blockCount * Q4_1BlockBytes, 64);
+        try
+        {
+            NativeMemory.Clear((void*)ptr, blockCount * Q4_1BlockBytes);
+            byte* b0 = (byte*)ptr;
+            byte* b1 = (byte*)ptr + Q4_1BlockBytes;
+
+            *(Half*)b0 = (Half)1.0f;  *(Half*)(b0 + 2) = (Half)0.0f;  b0[4] = 0x03;  // lo = 3
+            *(Half*)b1 = (Half)2.0f;  *(Half*)(b1 + 2) = (Half)1.0f;  b1[4] = 0x05;  // lo = 5
+
+            float[] dest = new float[blockCount * Q8_0GroupSize];
+            Dequantize.ToFloat32(ptr, blockCount * Q8_0GroupSize, QuantizationType.Q4_1, dest);
+
+            Assert.Equal(3.0f, dest[0], 1e-4f);   // 1.0 * 3 + 0
+            Assert.Equal(11.0f, dest[32], 1e-4f); // 2.0 * 5 + 1
+        }
+        finally
+        {
+            NativeMemory.AlignedFree((void*)ptr);
+        }
+    }
+
+    [Fact]
+    public void Q4_1_RowByteSize_Matches()
+    {
+        Assert.Equal(20L, Dequantize.RowByteSize(32, QuantizationType.Q4_1));
+        Assert.Equal(640L, Dequantize.RowByteSize(1024, QuantizationType.Q4_1));
+    }
+
+    [Fact]
+    public void Q4_1_NonAlignedCount_Throws()
+    {
+        float[] dest = new float[40];
+        Assert.Throws<ArgumentException>(() =>
+            Dequantize.ToFloat32(nint.Zero, 40, QuantizationType.Q4_1, dest));
+    }
+
+    // ──────────────────── Q5_1 ────────────────────
+
+    [Fact]
+    public void Q5_1_SingleBlock_HandCalculated()
+    {
+        // Layout (24 bytes / 32 elements): d(Half@0), m(Half@2), qh[4]@4, qs[16]@8.
+        // value = d * ((qh_bit << 4) | nibble) + m. Element j takes qh bit j,
+        // element j + 16 takes qh bit j + 16 (llama.cpp dequantize_row_q5_1).
+        // d = 0.25, m = 1. qs[0] = 0x21 → lo = 1, hi = 2. qh bit 0 set, bit 16 set.
+        //   dest[0]  = 0.25 * (16 | 1) + 1 = 0.25 * 17 + 1 = 5.25
+        //   dest[16] = 0.25 * (16 | 2) + 1 = 0.25 * 18 + 1 = 5.5
+        // qs[1] = 0x21 with no qh bits set:
+        //   dest[1]  = 0.25 * 1 + 1 = 1.25
+        //   dest[17] = 0.25 * 2 + 1 = 1.5
+        uint qh = (1u << 0) | (1u << 16);
+        nint ptr = AllocQ5_1Block((Half)0.25f, (Half)1.0f, qh, j => j <= 1 ? (byte)0x21 : (byte)0x00);
+        try
+        {
+            float[] dest = new float[Q5_0GroupSize];
+            Dequantize.ToFloat32(ptr, Q5_0GroupSize, QuantizationType.Q5_1, dest);
+
+            Assert.Equal(5.25f, dest[0], 1e-4f);
+            Assert.Equal(5.5f, dest[16], 1e-4f);
+            Assert.Equal(1.25f, dest[1], 1e-4f);
+            Assert.Equal(1.5f, dest[17], 1e-4f);
+        }
+        finally
+        {
+            NativeMemory.AlignedFree((void*)ptr);
+        }
+    }
+
+    [Fact]
+    public void Q5_1_AllBitsSet_GivesMaxCode()
+    {
+        // Every nibble 0xF and every high bit set → code 31 everywhere.
+        nint ptr = AllocQ5_1Block((Half)1.0f, (Half)0.0f, 0xFFFFFFFFu, _ => 0xFF);
+        try
+        {
+            float[] dest = new float[Q5_0GroupSize];
+            Dequantize.ToFloat32(ptr, Q5_0GroupSize, QuantizationType.Q5_1, dest);
+            for (int i = 0; i < Q5_0GroupSize; i++)
+                Assert.Equal(31.0f, dest[i], 1e-4f);
+        }
+        finally
+        {
+            NativeMemory.AlignedFree((void*)ptr);
+        }
+    }
+
+    [Fact]
+    public void Q5_1_TwoBlocks_StrideCorrect()
+    {
+        // Catches a wrong block stride (e.g. 22 bytes — the Q5_0 size — instead of 24).
+        const int blockCount = 2;
+        nint ptr = (nint)NativeMemory.AlignedAlloc(blockCount * Q5_1BlockBytes, 64);
+        try
+        {
+            NativeMemory.Clear((void*)ptr, blockCount * Q5_1BlockBytes);
+            byte* b0 = (byte*)ptr;
+            byte* b1 = (byte*)ptr + Q5_1BlockBytes;
+
+            *(Half*)b0 = (Half)1.0f;  *(Half*)(b0 + 2) = (Half)0.0f;  *(uint*)(b0 + 4) = 0u;  b0[8] = 0x07;
+            *(Half*)b1 = (Half)2.0f;  *(Half*)(b1 + 2) = (Half)1.0f;  *(uint*)(b1 + 4) = 1u;  b1[8] = 0x02;
+
+            float[] dest = new float[blockCount * Q5_0GroupSize];
+            Dequantize.ToFloat32(ptr, blockCount * Q5_0GroupSize, QuantizationType.Q5_1, dest);
+
+            Assert.Equal(7.0f, dest[0], 1e-4f);    // 1.0 * 7 + 0
+            Assert.Equal(37.0f, dest[32], 1e-4f);  // 2.0 * (16 | 2) + 1
+        }
+        finally
+        {
+            NativeMemory.AlignedFree((void*)ptr);
+        }
+    }
+
+    [Fact]
+    public void Q5_1_RowByteSize_Matches()
+    {
+        Assert.Equal(24L, Dequantize.RowByteSize(32, QuantizationType.Q5_1));
+        Assert.Equal(768L, Dequantize.RowByteSize(1024, QuantizationType.Q5_1));
+    }
+
+    [Fact]
+    public void Q5_1_NonAlignedCount_Throws()
+    {
+        float[] dest = new float[40];
+        Assert.Throws<ArgumentException>(() =>
+            Dequantize.ToFloat32(nint.Zero, 40, QuantizationType.Q5_1, dest));
+    }
+
     // ──────────────────── Helpers ────────────────────
 
+    private static nint AllocQ4_1Block(Half d, Half m, Func<int, byte> fillQs)
+    {
+        nint ptr = (nint)NativeMemory.AlignedAlloc(Q4_1BlockBytes, 32);
+        byte* p = (byte*)ptr;
+        *(Half*)p = d;
+        *(Half*)(p + 2) = m;
+        for (int i = 0; i < 16; i++)
+            (p + 4)[i] = fillQs(i);
+        return ptr;
+    }
+
+    private static nint AllocQ5_1Block(Half d, Half m, uint qh, Func<int, byte> fillQs)
+    {
+        nint ptr = (nint)NativeMemory.AlignedAlloc(Q5_1BlockBytes, 32);
+        byte* p = (byte*)ptr;
+        *(Half*)p = d;
+        *(Half*)(p + 2) = m;
+        *(uint*)(p + 4) = qh;
+        for (int i = 0; i < 16; i++)
+            (p + 8)[i] = fillQs(i);
+        return ptr;
+    }
+
     private static nint AllocQ8_0Block(Half scale, Func<int, sbyte> fillQs)
     {
         nint ptr = (nint)NativeMemory.AlignedAlloc(Q8_0BlockBytes, 32);
diff --git a/tests/DotLLM.Tests.Unit/Cuda/PtxTargetTests.cs b/tests/DotLLM.Tests.Unit/Cuda/PtxTargetTests.cs
new file mode 100644
index 00000000..67276999
--- /dev/null
+++ b/tests/DotLLM.Tests.Unit/Cuda/PtxTargetTests.cs
@@ -0,0 +1,50 @@
+using System.Text.RegularExpressions;
+using Xunit;
+
+namespace DotLLM.Tests.Unit.Cuda;
+
+/// <summary>
+/// Guards the compatibility baseline of the checked-in PTX. These files are build
+/// artifacts, so it is easy to regenerate one with a newer toolkit's default
+/// architecture and silently drop support for older GPUs and drivers — which has
+/// happened before (a CUDA 13.1 regeneration shipped <c>.target sm_75</c>).
+/// </summary>
+/// <remarks>
+/// The baseline is <c>compute_61</c>, matching <c>native/build.ps1</c> /
+/// <c>native/build.sh</c>. PTX is forward-compatible, so sm_61 PTX runs on every
+/// GPU from Pascal onward; sm_75 PTX does not load on Pascal at all.
+/// </remarks>
+public sealed class PtxTargetTests
+{
+    [Fact]
+    public void CheckedInPtx_TargetsBaselineArchitecture()
+    {
+        string ptxDir = FindPtxDir();
+        string[] files = Directory.GetFiles(ptxDir, "*.ptx");
+        Assert.NotEmpty(files);
+
+        foreach (string file in files)
+        {
+            string text = File.ReadAllText(file);
+            Match target = Regex.Match(text, @"^\.target\s+(\S+)", RegexOptions.Multiline);
+            Assert.True(target.Success, $"{Path.GetFileName(file)} declares no .target directive.");
+            Assert.True(
+                target.Groups[1].Value == "sm_61",
+                $"{Path.GetFileName(file)} targets '{target.Groups[1].Value}', not the sm_61 baseline. " +
+                "Regenerate with native/build.ps1 (or build.sh), which pins -arch=compute_61.");
+        }
+    }
+
+    private static string FindPtxDir()
+    {
+        var dir = new DirectoryInfo(AppContext.BaseDirectory);
+        while (dir != null)
+        {
+            string candidate = Path.Combine(dir.FullName, "native", "ptx");
+            if (Directory.Exists(candidate))
+                return candidate;
+            dir = dir.Parent;
+        }
+        throw new DirectoryNotFoundException("Could not locate native/ptx from the test output directory.");
+    }
+}