patches/0081-tcq-q6-asym-kv.patch

diff --git a/llama.cpp/BUILD.mk b/llama.cpp/BUILD.mk
index b678d3a3b..17f30f784 100644
--- a/llama.cpp/BUILD.mk
+++ b/llama.cpp/BUILD.mk
@@ -421,7 +421,7 @@ $(UI_CPP_GEN) $(UI_H_GEN) &: $(UI_EMBED_TOOL) $(UI_ASSETS_FILES)
 # Tool source files
 TOOL_QUANTIZE_SRCS := llama.cpp/tools/quantize/quantize.cpp
 TOOL_IMATRIX_SRCS := llama.cpp/tools/imatrix/imatrix.cpp
-TOOL_PERPLEXITY_SRCS := llama.cpp/tools/perplexity/perplexity.cpp
+TOOL_PERPLEXITY_SRCS := llama.cpp/tools/perplexity/perplexity.cpp llama.cpp/tools/perplexity/main.cpp
 TOOL_BENCH_SRCS := llama.cpp/tools/llama-bench/llama-bench.cpp
 
 TOOL_SERVER_SRCS := \
@@ -622,8 +622,12 @@ o/$(MODE)/llama.cpp/imatrix/imatrix: \
 
 o/$(MODE)/llama.cpp/perplexity/perplexity: \
 	$(TOOL_PERPLEXITY_OBJS) \
+	$(TOOL_LLAMAFILE_OBJS) \
+	$(HTTPLIB_OBJS) \
 	$$(TINYBLAS_CPU_OBJS) \
 	o/$(MODE)/llama.cpp/llama.cpp.a
+	@mkdir -p $(dir $@)
+	$(LINK.o) $(TOOL_PERPLEXITY_OBJS) $(TOOL_LLAMAFILE_OBJS) $(HTTPLIB_OBJS) $(TINYBLAS_CPU_OBJS) o/$(MODE)/llama.cpp/llama.cpp.a $(LOADLIBES) $(LDLIBS) -o $@
 
 o/$(MODE)/llama.cpp/llama-bench/llama-bench: \
 	$(TOOL_BENCH_OBJS) \
diff --git a/llama.cpp/common/arg.cpp b/llama.cpp/common/arg.cpp
index a5e6c9182..cfd3b293c 100644
--- a/llama.cpp/common/arg.cpp
+++ b/llama.cpp/common/arg.cpp
@@ -412,11 +412,17 @@ const std::vector<ggml_type> kv_cache_types = {
     GGML_TYPE_IQ4_NL,
     GGML_TYPE_Q5_0,
     GGML_TYPE_Q5_1,
+    // opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port) — -ctk/-ctv q6_0. See docs/features/poly_kv.md.
+    GGML_TYPE_Q6_0,
     // opencoti-hook: turboquant — see docs/features/poly_kv.md (M6-S2): -ctk/-ctv turbo2|turbo3|turbo4|turbo8
     GGML_TYPE_TURBO2_0,
     GGML_TYPE_TURBO3_0,
     GGML_TYPE_TURBO4_0,
     GGML_TYPE_TURBO8_0,
+    // opencoti-hook: TCQ (F5 M6 C5) — Trellis-Coded Quantization KV tiers (buun-aligned names):
+    // -ctk/-ctv turbo3_tcq|turbo2_tcq. See docs/evaluations/context.md §5.
+    GGML_TYPE_TURBO3_TCQ,
+    GGML_TYPE_TURBO2_TCQ,
 };
 
 static ggml_type kv_cache_type_from_str(const std::string & s) {
diff --git a/llama.cpp/ggml/include/ggml.h b/llama.cpp/ggml/include/ggml.h
index 3ab9e7775..3c72dd4b9 100644
--- a/llama.cpp/ggml/include/ggml.h
+++ b/llama.cpp/ggml/include/ggml.h
@@ -437,7 +437,18 @@ extern "C" {
         GGML_TYPE_TQ4_1S  = 46, // TurboQuant 4-bit weight: WHT-rotated Lloyd-Max, block=32
         // opencoti-hook: turboquant tier family (M6-S2) — 8-bit near-lossless KV rung
         GGML_TYPE_TURBO8_0 = 47, // TurboQuant 8-bit KV: WHT + uniform int8 (SCALE=256)
-        GGML_TYPE_COUNT   = 48,
+        // opencoti-hook: TCQ (F5 M6 C5) — Trellis-Coded Quantization KV tiers, ported from buun
+        // (github.com/spiritbuun/buun-llama-cpp). buun uses enum 45/46, but those are our
+        // TQ3_1S/TQ4_1S, so we take the free slots 48/49. KV cache types are string-selected at
+        // runtime (-ctk/-ctv) and NEVER GGUF-serialized, so the type NAME ("turbo3_tcq"/
+        // "turbo2_tcq") — not the enum value — is what must stay buun-compatible. See context.md §5.
+        GGML_TYPE_TURBO3_TCQ = 48, // TCQ 3-bit KV: FWHT + Viterbi (512-state, k=3, L=9), O(1) decode, 3.25 bpv
+        GGML_TYPE_TURBO2_TCQ = 49, // TCQ 2-bit KV: FWHT + Viterbi (256-state, k=2, L=8), O(1) decode, 2.25 bpv
+        // opencoti-hook: q6_0 KV — see docs/evaluations/context.md §5. Ported from Anbeeld/beellama.cpp
+        // (ik_llama lineage); the anbeeld.com KV-quant benchmark's recommended high-end V tier (q8_0/q6_0).
+        // KV-only (never GGUF-serialized); signed 6-bit, QK6_0=32, 6.5 bpv.
+        GGML_TYPE_Q6_0    = 50,
+        GGML_TYPE_COUNT   = 51,
     };
 
     // precision
diff --git a/llama.cpp/ggml/src/ggml-common.h b/llama.cpp/ggml/src/ggml-common.h
index 82d416e81..551fac0be 100644
--- a/llama.cpp/ggml/src/ggml-common.h
+++ b/llama.cpp/ggml/src/ggml-common.h
@@ -245,6 +245,17 @@ typedef struct {
 } block_q8_0;
 static_assert(sizeof(block_q8_0) == sizeof(ggml_half) + QK8_0, "wrong q8_0 block size/padding");
 
+// opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port, ik_llama lineage) — see docs/evaluations/context.md §5.
+// Signed 6-bit: qs[16] holds the low 4 bits of all 32 values; qh[8] holds the high 2 bits packed 4-per-byte.
+// 2 + 8 + 16 = 26 bytes / 32 values = 6.5 bpv. KV-only (never GGUF-serialized).
+#define QK6_0 32
+typedef struct {
+    ggml_half d;           // delta
+    uint8_t qh[QK6_0/4];   // high 2 bits of quants (packed 4 values/byte)
+    uint8_t qs[QK6_0/2];   // low 4 bits / quants
+} block_q6_0;
+static_assert(sizeof(block_q6_0) == sizeof(ggml_half) + QK6_0/4 + QK6_0/2, "wrong q6_0 block size/padding");
+
 #define QK8_1 32
 typedef struct {
     GGML_EXTENSION union {
@@ -531,6 +542,33 @@ typedef struct {
 } block_turbo8_0;                       // 130 bytes total
 static_assert(sizeof(block_turbo8_0) == sizeof(ggml_half) + QK_TURBO8, "wrong turbo8_0 block size/padding");
 
+// opencoti-hook: TCQ (F5 M6 C5) — Trellis-Coded Quantization KV, ported from buun
+// (github.com/spiritbuun/buun-llama-cpp). Real encode (FWHT + Viterbi) is the CUDA SET_ROWS
+// kernel; real decode (O(1) sliding-window codebook lookup) is in the FA-VEC kernel. CPU
+// to_float/from_float are deliberate stubs (norm-only / zero-fill). block == rotation group == 128.
+//
+// TurboQuant 3-bit TCQ: Trellis-Coded Quantization (right-shift bitshift trellis, k=3, L=9)
+// One block = one 128-element rotation group. Bitstream: 6 zero-prefix + 128×3-bit = 390 bits = 49 bytes.
+// Decode: state_t = read_9_bits(qs, t*3), recon_t = codebook[state_t] * norm. 3.25 bpv → 5.0× vs fp16.
+#define QK_TURBO3_TCQ 128
+typedef struct {
+    ggml_half  norm;                    //  2 bytes: corrected group L2 norm
+    uint8_t    qs[49];                  // 49 bytes: 390-bit trellis bitstream (2 padding bits)
+    uint8_t    pad;                     //  1 byte:  alignment padding
+} block_turbo3_tcq;                     // 52 bytes total for 128 values (3.25 bpv)
+static_assert(sizeof(block_turbo3_tcq) == sizeof(ggml_half) + 50, "wrong turbo3_tcq block size/padding");
+
+// TurboQuant 2-bit TCQ: Trellis-Coded Quantization (right-shift bitshift trellis, k=2, L=8)
+// One block = one 128-element rotation group. Bitstream: 6 prefix + 128×2-bit = 262 bits = 33 bytes.
+// Decode: state_t = read_8_bits(qs, t*2), recon_t = codebook[state_t] * norm. 2.25 bpv → 7.1× vs fp16.
+#define QK_TURBO2_TCQ 128
+typedef struct {
+    ggml_half  norm;                    //  2 bytes: corrected group L2 norm
+    uint8_t    qs[33];                  // 33 bytes: 262-bit trellis bitstream (2 padding bits)
+    uint8_t    pad;                     //  1 byte:  alignment padding
+} block_turbo2_tcq;                     // 36 bytes total for 128 values (2.25 bpv)
+static_assert(sizeof(block_turbo2_tcq) == sizeof(ggml_half) + 34, "wrong turbo2_tcq block size/padding");
+
 // TQ3_1S: WHT-rotated 3-bit weight quantization (8-level Lloyd-Max for N(0,1))
 // Block size 32, dual half-block scales (d0 for [0..15], d1 for [16..31])
 // Per block: d0(fp16) + d1(fp16) + 3-bit indices packed (12 bytes) = 16 bytes per 32 values
diff --git a/llama.cpp/ggml/src/ggml-cpu/ggml-cpu.c b/llama.cpp/ggml/src/ggml-cpu/ggml-cpu.c
index aad25a6ce..468f762db 100644
--- a/llama.cpp/ggml/src/ggml-cpu/ggml-cpu.c
+++ b/llama.cpp/ggml/src/ggml-cpu/ggml-cpu.c
@@ -268,6 +268,11 @@ static const struct ggml_type_traits_cpu type_traits_cpu[GGML_TYPE_COUNT] = {
         .vec_dot_type             = GGML_TYPE_Q8_1,
         .nrows                    = 1,
     },
+    // opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port) — KV-only scalar 6-bit; from_float via the
+    // ref (no CPU vec_dot — dequant-on-lift converts q6_0→f16 before any matmul). See docs/features/poly_kv.md.
+    [GGML_TYPE_Q6_0] = {
+        .from_float               = (ggml_from_float_t) quantize_row_q6_0_ref,
+    },
     [GGML_TYPE_Q8_0] = {
         .from_float               = quantize_row_q8_0,
         .vec_dot                  = ggml_vec_dot_q8_0_q8_0,
@@ -433,6 +438,15 @@ static const struct ggml_type_traits_cpu type_traits_cpu[GGML_TYPE_COUNT] = {
     [GGML_TYPE_TURBO8_0] = {
         .from_float               = (ggml_from_float_t) quantize_row_turbo8_0_ref,
     },
+    // opencoti-hook: TCQ (F5 M6 C5) — Trellis-Coded Quantization KV. CPU from_float is a norm-only
+    // stub (real Viterbi encode is the CUDA set_rows kernel); registered so the dequant-on-lift /
+    // POSITION_WINDOW host-pinned tail set_rows resolves a from_float for the type.
+    [GGML_TYPE_TURBO3_TCQ] = {
+        .from_float               = (ggml_from_float_t) quantize_row_turbo3_tcq_ref,
+    },
+    [GGML_TYPE_TURBO2_TCQ] = {
+        .from_float               = (ggml_from_float_t) quantize_row_turbo2_tcq_ref,
+    },
 };
 
 const struct ggml_type_traits_cpu * ggml_get_type_traits_cpu(enum ggml_type type) {
diff --git a/llama.cpp/ggml/src/ggml-cpu/quants.h b/llama.cpp/ggml/src/ggml-cpu/quants.h
index 950b2a916..922db2508 100644
--- a/llama.cpp/ggml/src/ggml-cpu/quants.h
+++ b/llama.cpp/ggml/src/ggml-cpu/quants.h
@@ -24,6 +24,14 @@ void quantize_row_turbo3_0_ref(const float * GGML_RESTRICT x, block_turbo3_0 * G
 void quantize_row_turbo4_0_ref(const float * GGML_RESTRICT x, block_turbo4_0 * GGML_RESTRICT y, int64_t k);
 void quantize_row_turbo2_0_ref(const float * GGML_RESTRICT x, block_turbo2_0 * GGML_RESTRICT y, int64_t k);
 void quantize_row_turbo8_0_ref(const float * GGML_RESTRICT x, block_turbo8_0 * GGML_RESTRICT y, int64_t k);
+// opencoti-hook: TCQ (F5 M6 C5) — CPU from_float refs (norm-only stubs) must also be declared here:
+// ggml-cpu.c includes this CPU-local quants.h, NOT the global ggml-quants.h, for the from_float table.
+void quantize_row_turbo3_tcq_ref(const float * GGML_RESTRICT x, block_turbo3_tcq * GGML_RESTRICT y, int64_t k);
+void quantize_row_turbo2_tcq_ref(const float * GGML_RESTRICT x, block_turbo2_tcq * GGML_RESTRICT y, int64_t k);
+// opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port) — same CPU-local-header gotcha: ggml-cpu.c's q6_0
+// from_float = quantize_row_q6_0_ref resolves against THIS header, not the global ggml-quants.h, so the
+// q6_0 ref must be declared here too (defined in ggml-quants.c).
+void quantize_row_q6_0_ref(const float * GGML_RESTRICT x, block_q6_0 * GGML_RESTRICT y, int64_t k);
 
 void quantize_row_mxfp4(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
 void quantize_row_nvfp4(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
diff --git a/llama.cpp/ggml/src/ggml-cuda/common.cuh b/llama.cpp/ggml/src/ggml-cuda/common.cuh
index c9275517f..1fdf9d02a 100644
--- a/llama.cpp/ggml/src/ggml-cuda/common.cuh
+++ b/llama.cpp/ggml/src/ggml-cuda/common.cuh
@@ -21,6 +21,19 @@
 #endif
 #include "ggml-common.h"
 
+// opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port). The host ggml-common.h defines QK6_0 (32) but
+// not the QR/QI dequant-reduction macros (q5_0's QR5_0/QI5_0 live there; q6_0's were never added
+// host-side, and host files are frozen for this port). Define them CUDA-side here, right after
+// ggml-common.h (where QK6_0 lands), so convert.cu's dequantize_block_cont_cuda<QK6_0,QR6_0,...> and
+// fattn-common.cuh's K reader (QI6_0) resolve. q6_0 packs 2 values per qs byte ⇒ QR6_0 = 2 (mirrors
+// QR5_0); QI6_0 = QK6_0/(4*QR6_0) = 4 (int32 lanes of qs per block, used by the FA-VEC iqs index).
+#ifndef QR6_0
+#define QR6_0 2
+#endif
+#ifndef QI6_0
+#define QI6_0 (QK6_0 / (4 * QR6_0))
+#endif
+
 #include <array>
 #include <algorithm>
 #include <cassert>
diff --git a/llama.cpp/ggml/src/ggml-cuda/convert.cu b/llama.cpp/ggml/src/ggml-cuda/convert.cu
index 0d27c1672..6ee9aa92b 100644
--- a/llama.cpp/ggml/src/ggml-cuda/convert.cu
+++ b/llama.cpp/ggml/src/ggml-cuda/convert.cu
@@ -720,6 +720,8 @@ to_fp16_cuda_t ggml_get_to_fp16_cuda(ggml_type type) {
             return dequantize_row_q4_1_cuda;
         case GGML_TYPE_Q5_0:
             return dequantize_block_cont_cuda<QK5_0, QR5_0, dequantize_q5_0>;
+        case GGML_TYPE_Q6_0: // opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port)
+            return dequantize_block_cont_cuda<QK6_0, QR6_0, dequantize_q6_0>;
         case GGML_TYPE_Q5_1:
             return dequantize_block_cont_cuda<QK5_1, QR5_1, dequantize_q5_1>;
         case GGML_TYPE_Q8_0:
@@ -789,6 +791,8 @@ to_fp32_cuda_t ggml_get_to_fp32_cuda(ggml_type type) {
             return dequantize_row_q4_1_cuda;
         case GGML_TYPE_Q5_0:
             return dequantize_block_cont_cuda<QK5_0, QR5_0, dequantize_q5_0>;
+        case GGML_TYPE_Q6_0: // opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port)
+            return dequantize_block_cont_cuda<QK6_0, QR6_0, dequantize_q6_0>;
         case GGML_TYPE_Q5_1:
             return dequantize_block_cont_cuda<QK5_1, QR5_1, dequantize_q5_1>;
         case GGML_TYPE_Q8_0:
@@ -848,6 +852,8 @@ to_fp16_nc_cuda_t ggml_get_to_fp16_nc_cuda(ggml_type type) {
             return dequantize_block_cuda<QK4_1, QR4_1, dequantize_q4_1>;
         case GGML_TYPE_Q5_0:
             return dequantize_block_cuda<QK5_0, QR5_0, dequantize_q5_0>;
+        case GGML_TYPE_Q6_0: // opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port)
+            return dequantize_block_cuda<QK6_0, QR6_0, dequantize_q6_0>;
         case GGML_TYPE_Q5_1:
             return dequantize_block_cuda<QK5_1, QR5_1, dequantize_q5_1>;
         case GGML_TYPE_Q8_0:
@@ -880,6 +886,8 @@ to_bf16_nc_cuda_t ggml_get_to_bf16_nc_cuda(ggml_type type) {
             return dequantize_block_cuda<QK4_1, QR4_1, dequantize_q4_1>;
         case GGML_TYPE_Q5_0:
             return dequantize_block_cuda<QK5_0, QR5_0, dequantize_q5_0>;
+        case GGML_TYPE_Q6_0: // opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port)
+            return dequantize_block_cuda<QK6_0, QR6_0, dequantize_q6_0>;
         case GGML_TYPE_Q5_1:
             return dequantize_block_cuda<QK5_1, QR5_1, dequantize_q5_1>;
         case GGML_TYPE_Q8_0:
@@ -903,6 +911,8 @@ to_fp32_nc_cuda_t ggml_get_to_fp32_nc_cuda(ggml_type type) {
             return dequantize_block_cuda<QK4_1, QR4_1, dequantize_q4_1>;
         case GGML_TYPE_Q5_0:
             return dequantize_block_cuda<QK5_0, QR5_0, dequantize_q5_0>;
+        case GGML_TYPE_Q6_0: // opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port)
+            return dequantize_block_cuda<QK6_0, QR6_0, dequantize_q6_0>;
         case GGML_TYPE_Q5_1:
             return dequantize_block_cuda<QK5_1, QR5_1, dequantize_q5_1>;
         case GGML_TYPE_Q8_0:
diff --git a/llama.cpp/ggml/src/ggml-cuda/cpy-utils.cuh b/llama.cpp/ggml/src/ggml-cuda/cpy-utils.cuh
index 7697c292d..d82afeb18 100644
--- a/llama.cpp/ggml/src/ggml-cuda/cpy-utils.cuh
+++ b/llama.cpp/ggml/src/ggml-cuda/cpy-utils.cuh
@@ -103,6 +103,43 @@ static __device__ void quantize_f32_q5_0_block(const float * __restrict__ x, blo
     memcpy(y->qh, &qh, sizeof(qh));
 }
 
+// opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port). f32->q6_0 block pack, modeled on
+// quantize_f32_q5_0_block above (our amax/vmax scan + per-half loop idiom). Pack math ported verbatim
+// from ik_llama.cpp cpy-utils.cuh:210-241 (quantize_f32_q6_0_block): d = vmax/-32 (NO sumqx/sumq2 refit
+// for q6_0 — ik only refits iq4_nl); 6-bit codes via min(63, (int8_t)(x*id + 32.5)); qs holds low
+// nibbles, qh[8] holds the high 2 bits packed 4 values/byte (h = (xi0>>4)|((xi1>>4)<<2), then
+// qh[j%8] |= h << 4*(j/8)). Produces bytes byte-identical to the host quantize_row_q6_0_ref.
+static __device__ void quantize_f32_q6_0_block(const float * __restrict__ x, block_q6_0 * __restrict__ y) {
+    float amax = 0.0f;
+    float vmax = 0.0f;
+
+    for (int j = 0; j < QK6_0; ++j) {
+        const float v = x[j];
+        if (amax < fabsf(v)) {
+            amax = fabsf(v);
+            vmax = v;
+        }
+    }
+
+    const float d  = vmax / -32;
+    const float id = d ? 1.0f/d : 0.0f;
+
+    y->d = d;
+    memset(y->qh, 0, QK6_0/4);
+
+    for (int j = 0; j < QK6_0/2; ++j) {
+        const float x0 = x[0       + j]*id;
+        const float x1 = x[QK6_0/2 + j]*id;
+
+        const uint8_t xi0 = min(63, (int8_t)(x0 + 32.5f));
+        const uint8_t xi1 = min(63, (int8_t)(x1 + 32.5f));
+
+        y->qs[j]  = (xi0 & 0xf) | ((xi1 & 0xf) << 4);
+        const uint8_t h = (xi0 >> 4) | ((xi1 >> 4) << 2);
+        y->qh[j % (QK6_0/4)] |= (h << 4*(j / (QK6_0/4)));
+    }
+}
+
 static __device__ void quantize_f32_q5_1_block(const float * __restrict__ x, block_q5_1 * __restrict__ y) {
     float min = x[0];
     float max = x[0];
@@ -203,6 +240,11 @@ static __device__ void cpy_blck_f32_q5_1(const char * cxi, char * cdsti) {
     quantize_f32_q5_1_block((const float *)cxi, (block_q5_1 *)cdsti);
 }
 
+// opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port). cpy.cu wrapper, mirrors cpy_blck_f32_q5_0.
+static __device__ void cpy_blck_f32_q6_0(const char * cxi, char * cdsti) {
+    quantize_f32_q6_0_block((const float *)cxi, (block_q6_0 *)cdsti);
+}
+
 static __device__ void cpy_blck_f32_q8_0(const char * cxi, char * cdsti) {
     quantize_f32_q8_0_block((const float *)cxi, (block_q8_0 *)cdsti);
 }
diff --git a/llama.cpp/ggml/src/ggml-cuda/cpy.cu b/llama.cpp/ggml/src/ggml-cuda/cpy.cu
index 5e01613fd..e769b7bf7 100644
--- a/llama.cpp/ggml/src/ggml-cuda/cpy.cu
+++ b/llama.cpp/ggml/src/ggml-cuda/cpy.cu
@@ -344,6 +344,35 @@ static void ggml_cpy_q5_0_f32_cuda(
         ne10, ne11, ne12, nb10, nb11, nb12, nb13);
 }
 
+// opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port). f32->q6_0 / q6_0->f32 cpy launchers, mirror
+// ggml_cpy_f32_q5_0_cuda / ggml_cpy_q5_0_f32_cuda above with QK6_0 block size and the q6_0
+// pack/dequant block functions (cpy_blck_f32_q6_0 / dequantize_q6_0).
+static void ggml_cpy_f32_q6_0_cuda(
+    const char * cx, char * cdst, const int64_t ne,
+    const int64_t ne00, const int64_t ne01, const int64_t ne02, const int64_t nb00, const int64_t nb01, const int64_t nb02,
+    const int64_t nb03, const int64_t ne10, const int64_t ne11, const int64_t ne12, const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13, cudaStream_t stream) {
+
+    GGML_ASSERT(ne % QK6_0 == 0);
+    const int64_t num_blocks = ne / QK6_0;
+    GGML_ASSERT(num_blocks < UINT_MAX);
+    cpy_f32_q<cpy_blck_f32_q6_0, QK6_0><<<num_blocks, 1, 0, stream>>>
+        (cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13);
+}
+
+static void ggml_cpy_q6_0_f32_cuda(
+    const char * cx, char * cdst, const int64_t ne,
+    const int64_t ne00, const int64_t ne01, const int64_t ne02,
+    const int64_t nb00, const int64_t nb01, const int64_t nb02,
+    const int64_t nb03, const int64_t ne10, const int64_t ne11, const int64_t ne12,
+    const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13,
+    cudaStream_t stream) {
+    const int64_t num_blocks = ne;
+    GGML_ASSERT(num_blocks < UINT_MAX);
+    cpy_q_f32<cpy_blck_q_f32<dequantize_q6_0, QK6_0>, QK6_0><<<num_blocks, 1, 0, stream>>>(
+        cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03,
+        ne10, ne11, ne12, nb10, nb11, nb12, nb13);
+}
+
 static void ggml_cpy_f32_q5_1_cuda(
     const char * cx, char * cdst, const int64_t ne,
     const int64_t ne00, const int64_t ne01, const int64_t ne02, const int64_t nb00, const int64_t nb01, const int64_t nb02,
@@ -477,6 +506,13 @@ void ggml_cuda_cpy(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, gg
     } else if (src0->type == GGML_TYPE_Q5_0 && src1->type == GGML_TYPE_F32) {
         ggml_cpy_q5_0_f32_cuda
                 (src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
+    // opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port) — f32<->q6_0 cpy dispatch, mirrors q5_0.
+    } else if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_Q6_0) {
+        ggml_cpy_f32_q6_0_cuda
+                (src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
+    } else if (src0->type == GGML_TYPE_Q6_0 && src1->type == GGML_TYPE_F32) {
+        ggml_cpy_q6_0_f32_cuda
+                (src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream);
 #ifndef GGML_CUDA_NO_IQ_QUANTS
     } else if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_IQ4_NL) {
         ggml_cpy_f32_iq4_nl_cuda
diff --git a/llama.cpp/ggml/src/ggml-cuda/dequantize.cuh b/llama.cpp/ggml/src/ggml-cuda/dequantize.cuh
index 9ae1342fc..8e022111e 100644
--- a/llama.cpp/ggml/src/ggml-cuda/dequantize.cuh
+++ b/llama.cpp/ggml/src/ggml-cuda/dequantize.cuh
@@ -68,6 +68,25 @@ static __device__ __forceinline__ void dequantize_q5_0(const void * vx, const in
     v.y = (v.y - 16.0f) * d;
 }
 
+// opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port). Scalar block dequantizer (2 values/call),
+// modeled on dequantize_q5_0 above (our float2& idiom; our tree has no GGML_CUDA_F16 branch here).
+// Bit-unpack ported verbatim from ik_llama.cpp dequantize.cuh:107-123 (dequantize_q6_0):
+// h = qh[iqs%8] >> 4*(iqs/8); value iqs = (qs[iqs]&0xf) | ((h&0x3)<<4); value iqs+16 =
+// (qs[iqs]>>4) | ((h&0xc)<<2); both minus the q6_0 level offset 32. Used by convert.cu (to_fpXX)
+// and cpy.cu (q6_0->f32) via dequantize_block_cont_cuda<QK6_0,QR6_0,dequantize_q6_0>.
+static __device__ __forceinline__ void dequantize_q6_0(const void * vx, const int64_t ib, const int iqs, float2 & v){
+    const block_q6_0 * x = (const block_q6_0 *) vx;
+
+    const float d = x[ib].d;
+
+    const uint8_t h = x[ib].qh[iqs%8] >> 4*(iqs/8);
+    v.x = ((x[ib].qs[iqs] & 0xf) | ((h & 0x3) << 4));
+    v.y = ((x[ib].qs[iqs] >>  4) | ((h & 0xc) << 2));
+
+    v.x = (v.x - 32.0f) * d;
+    v.y = (v.y - 32.0f) * d;
+}
+
 static __device__ __forceinline__ void dequantize_q5_1(const void * vx, const int64_t ib, const int iqs, float2 & v){
     const block_q5_1 * x = (const block_q5_1 *) vx;
 
diff --git a/llama.cpp/ggml/src/ggml-cuda/fattn-common.cuh b/llama.cpp/ggml/src/ggml-cuda/fattn-common.cuh
index f61cd9d37..e5181d703 100644
--- a/llama.cpp/ggml/src/ggml-cuda/fattn-common.cuh
+++ b/llama.cpp/ggml/src/ggml-cuda/fattn-common.cuh
@@ -232,6 +232,54 @@ static __device__ __forceinline__ float vec_dot_fattn_vec_KQ_q5_0(
     return sum;
 }
 
+// opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port). K reader modeled on vec_dot_fattn_vec_KQ_q5_0
+// above (our <D,nthreads> float-return idiom), with the q6_0 bit-unpack math ported verbatim from
+// ik_llama.cpp fattn-common.cuh:290-331 (vec_dot_fattn_vec_KQ_q6_0). Differences from q5_0: q6_0 packs
+// the high 2 bits across qh[8] (4 values/byte) rather than 1 bit/value in a uint32 qh, the level
+// offset is 32 (not 16), and the Q_ds combine uses (32/QI8_1) instead of (16/QI8_1). q6_0 carries no
+// min term (symmetric, like q5_0) so the float path mirrors q5_0's single sumi*Q_ds.x - off*Q_ds.y.
+template<int D, int nthreads>
+static __device__ __forceinline__ float vec_dot_fattn_vec_KQ_q6_0(
+    const char * __restrict__ K_c, const void * __restrict__ Q_v, const int * __restrict__ Q_q8, const void * __restrict__ Q_ds_v) {
+
+    const block_q6_0 * K_q6_0 = (const block_q6_0 *) K_c;
+    GGML_UNUSED(Q_v);
+
+    float sum = 0.0f;
+
+#pragma unroll
+    for (int k_KQ_0 = 0; k_KQ_0 < int(D/sizeof(int)); k_KQ_0 += nthreads) {
+        const int k_KQ = k_KQ_0 + (nthreads == WARP_SIZE ? threadIdx.x : threadIdx.x % nthreads);
+
+        const int ib    = k_KQ /  QI8_1;
+        const int iqs4  = k_KQ %  QI6_0;       // 0...3
+        const int shift = k_KQ & (QI8_1/2);
+
+        // High 2 bits packed across qh[8] (4 values/byte); ik selects the qh int by iqs4%2 and shifts
+        // by 4*(iqs4/2) + shift/2 to land on this lane's 2-bit field, masked to 0x03 per byte.
+        int vh;
+        ggml_cuda_memcpy_1<sizeof(int), 2>(&vh, K_q6_0[ib].qh + sizeof(int)*(iqs4 % 2));
+        vh = (vh >> (4*(iqs4/2) + shift/2)) & 0x03030303;
+
+        // Low 4 bits: one int (4 values) of qs at lane iqs4, shifted by `shift` (0 or 4), masked to nibble.
+        int vl;
+        ggml_cuda_memcpy_1<sizeof(int), 2>(&vl, K_q6_0[ib].qs + sizeof(int)*iqs4);
+        vl = (vl >> shift) & 0x0F0F0F0F;
+
+        const int v = vl | (vh << 4);
+
+        const int u = Q_q8[k_KQ_0/nthreads];
+
+        const int sumi = ggml_cuda_dp4a(v, u, 0);
+
+        const float2 Q_ds = ((const float2 *) Q_ds_v)[k_KQ_0/nthreads];
+
+        sum += __half2float(K_q6_0[ib].d) * (sumi*Q_ds.x - (32/QI8_1)*Q_ds.y);
+    }
+
+    return sum;
+}
+
 template<int D, int nthreads>
 static __device__ __forceinline__ float vec_dot_fattn_vec_KQ_q5_1(
     const char * __restrict__ K_c, const void * __restrict__ Q_v, const int * __restrict__ Q_q8, const void * __restrict__ Q_ds_v) {
@@ -512,6 +560,66 @@ static __device__ __forceinline__ void dequantize_V_q5_0(const void * __restrict
     }
 }
 
+// opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port). V reader modeled on dequantize_V_q5_0 above
+// (our <T,ne> block-fill idiom: read ne low-nibbles into a packed int, OR in the high bits per lane,
+// subtract the level offset, scale by d). The high-bit unpack + offset(32) math is ported from
+// ik_llama.cpp fattn-common.cuh:566-588 (dequantize_1_q6_0), generalized to ne lanes: for output lane
+// l at absolute index idq+l, the high 2 bits come from qh[(idq+l)%8] >> 4*((idq+l)/8), then &0x03<<4.
+template <typename T, int ne>
+static __device__ __forceinline__ void dequantize_V_q6_0(const void * __restrict__ vx, void * __restrict__ dst, const int64_t i0) {
+    const block_q6_0 * x = (const block_q6_0 *) vx;
+
+    const int64_t ib    =  i0          /  QK6_0;
+    const int     idq   =  i0          %  QK6_0;
+    const int     iqs   =  i0          % (QK6_0/2);
+    const int     shift = (i0 % QK6_0) / (QK6_0/2);
+
+    int q;
+    static_assert(ne == 2 || ne == 4, "bad ne");
+    ggml_cuda_memcpy_1<ne, 2>(&q, x[ib].qs + iqs);
+    q >>= 4*shift;
+    q &= 0x0F0F0F0F;
+
+    {
+        // High 2 bits, ported verbatim from ik_llama.cpp dequantize_1_q6_0 (fattn-common.cuh:577-578)
+        // and evaluated per output lane l at value index (idq + l). For value V: byte = qh[V%(QK6_0/4)],
+        // 2-bit field at shift = 4*((V/(QK6_0/4))%2) + 2*shift. `shift` is constant across the ne lanes
+        // of one call (i0 is aligned within a nibble-half), so it equals (idq+l)/(QK6_0/2) for every l.
+        // Verified byte-exact against the host quantize_row_q6_0 pack for all 32 block positions.
+#pragma unroll
+        for (int l = 0; l < ne; ++l) {
+            const int j = idq + l;
+            const int h = (int) x[ib].qh[j % (QK6_0/4)] >> (4*((j / (QK6_0/4)) % 2) + 2*shift);
+            q |= (h & 0x03) << (8*l + 4);
+        }
+    }
+
+    q = __vsubss4(q, 0x20202020); // q6_0 level offset is 32 (0x20), vs q5_0's 16
+
+    const int8_t * q8 = (const int8_t *) &q;
+
+#ifdef FP16_AVAILABLE
+    if constexpr (std::is_same_v<T, half>) {
+        const half2 d = __half2half2(x[ib].d);
+
+#pragma unroll
+        for (int l0 = 0; l0 < ne; l0 += 2) {
+            ((half2 *) dst)[l0/2] = d * make_half2(q8[l0 + 0], q8[l0 + 1]);
+        }
+    } else
+#endif // FP16_AVAILABLE
+    if constexpr (std::is_same_v<T, float>) {
+        const float d = x[ib].d;
+
+#pragma unroll
+        for (int l = 0; l < ne; ++l) {
+            ((float *) dst)[l] = d * q8[l];
+        }
+    } else {
+        static_assert(std::is_same_v<T, void>, "bad type");
+    }
+}
+
 template <typename T, int ne>
 static __device__ __forceinline__ void dequantize_V_q5_1(const void * __restrict__ vx, void * __restrict__ dst, const int64_t i0) {
     const block_q5_1 * x = (const block_q5_1 *) vx;
@@ -599,6 +707,10 @@ static __device__ __forceinline__ void dequantize_V_q8_0(const void * __restrict
 // (centroid×norm, rotated space) for the fused FA-VEC path. Included here so get_vec_dot_KQ /
 // get_dequantize_V below can dispatch turbo2/3/4. ggml_cuda_mad + block_turbo* are already in scope.
 #include "fattn-turbo.cuh"
+// opencoti-hook: TCQ (trellis-coded quant) FA-VEC decode — F5 M6 C5 (#447/#500). Decode codebooks +
+// O(1) sliding-window trellis readers, ported verbatim from buun. Included after fattn-turbo.cuh so
+// block_turbo{2,3}_tcq / QK_TURBO{2,3}_TCQ / cpy-bytes helper are in scope. See UPSTREAM_SYNC.md.
+#include "fattn-tcq.cuh"
 
 template <ggml_type type_K, int D, int nthreads>
 constexpr __device__ vec_dot_KQ_t get_vec_dot_KQ() {
@@ -610,12 +722,18 @@ constexpr __device__ vec_dot_KQ_t get_vec_dot_KQ() {
         return vec_dot_fattn_vec_KQ_turbo3<D, nthreads>;
     } else if constexpr (type_K == GGML_TYPE_TURBO4_0) {
         return vec_dot_fattn_vec_KQ_turbo4<D, nthreads>;
+    } else if constexpr (type_K == GGML_TYPE_TURBO3_TCQ) { // opencoti-hook: TCQ FA-VEC decode (#447/#500)
+        return vec_dot_fattn_vec_KQ_turbo3_tcq<D, nthreads>;
+    } else if constexpr (type_K == GGML_TYPE_TURBO2_TCQ) {
+        return vec_dot_fattn_vec_KQ_turbo2_tcq<D, nthreads>;
     } else if constexpr (type_K == GGML_TYPE_Q4_0) {
         return vec_dot_fattn_vec_KQ_q4_0<D, nthreads>;
     } else if constexpr (type_K == GGML_TYPE_Q4_1) {
         return vec_dot_fattn_vec_KQ_q4_1<D, nthreads>;
     } else if constexpr (type_K == GGML_TYPE_Q5_0) {
         return vec_dot_fattn_vec_KQ_q5_0<D, nthreads>;
+    } else if constexpr (type_K == GGML_TYPE_Q6_0) { // opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port)
+        return vec_dot_fattn_vec_KQ_q6_0<D, nthreads>;
     } else if constexpr (type_K == GGML_TYPE_Q5_1) {
         return vec_dot_fattn_vec_KQ_q5_1<D, nthreads>;
     } else if constexpr (type_K == GGML_TYPE_Q8_0) {
@@ -638,12 +756,18 @@ constexpr __device__ dequantize_V_t get_dequantize_V() {
         return dequantize_V_turbo3<T, ne>;
     } else if constexpr (type_V == GGML_TYPE_TURBO4_0) {
         return dequantize_V_turbo4<T, ne>;
+    } else if constexpr (type_V == GGML_TYPE_TURBO3_TCQ) { // opencoti-hook: TCQ FA-VEC decode (#447/#500)
+        return dequantize_V_turbo3_tcq<T, ne>;
+    } else if constexpr (type_V == GGML_TYPE_TURBO2_TCQ) {
+        return dequantize_V_turbo2_tcq<T, ne>;
     } else if constexpr (type_V == GGML_TYPE_Q4_0) {
         return dequantize_V_q4_0<T, ne>;
     } else if constexpr (type_V == GGML_TYPE_Q4_1) {
         return dequantize_V_q4_1<T, ne>;
     } else if constexpr (type_V == GGML_TYPE_Q5_0) {
         return dequantize_V_q5_0<T, ne>;
+    } else if constexpr (type_V == GGML_TYPE_Q6_0) { // opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port)
+        return dequantize_V_q6_0<T, ne>;
     } else if constexpr (type_V == GGML_TYPE_Q5_1) {
         return dequantize_V_q5_1<T, ne>;
     } else if constexpr (type_V == GGML_TYPE_Q8_0) {
diff --git a/llama.cpp/ggml/src/ggml-cuda/fattn-tcq.cuh b/llama.cpp/ggml/src/ggml-cuda/fattn-tcq.cuh
new file mode 100644
index 000000000..a9541f575
--- /dev/null
+++ b/llama.cpp/ggml/src/ggml-cuda/fattn-tcq.cuh
@@ -0,0 +1,304 @@
+#pragma once
+// opencoti-hook: TCQ (trellis-coded quant) FA-VEC decode — F5 M6 C5 (#447/#500). See docs/protocols/UPSTREAM_SYNC.md.
+// Ported VERBATIM from buun-llama-cpp (github.com/spiritbuun/buun-llama-cpp @87c351d) fattn-common.cuh:
+// decode codebooks (cross-validated byte-identical to the encode codebooks in turbo-tcq-cuda.cuh),
+// context-adaptive decode alphas, and the O(1) sliding-window trellis readers (9-bit/512-state for
+// turbo3, 8-bit/256-state for turbo2). buun naming retained per project decision.
+// Included by fattn-common.cuh AFTER fattn-turbo.cuh, so these are already in scope:
+//   block_turbo{2,3}_tcq, QK_TURBO{2,3}_TCQ (=128), ggml_cuda_get_max_cpy_bytes(),
+//   V_DOT2_F32_F16_AVAILABLE, FP16_AVAILABLE, GGML_UNUSED.
+
+// --- 3-bit TCQ decode codebook (512-state), verbatim from buun fattn-common.cuh:60-125 ---
+static __constant__ float d_turbo3_tcq_codebook_fattn[512] = {
+    -0.14559399f, -0.09062801f, -0.054925077f, -0.03699251f, -0.006363985f, +0.026264573f, +0.067378916f, +0.121981815f,
+    -0.18648055f, -0.106522456f, -0.052047577f, -0.011695214f, +0.021953275f, +0.059698727f, +0.09831437f, +0.16083933f,
+    -0.16390342f, -0.12639847f, -0.09513180f, -0.05938352f, -0.028396897f, +0.005973862f, +0.049104784f, +0.11334257f,
+    -0.25952467f, -0.079778515f, -0.036024813f, +0.0003641268f, +0.031858794f, +0.073280424f, +0.11835553f, +0.19738495f,
+    -0.14218009f, -0.10224814f, -0.062498566f, -0.027066832f, +0.00393002f, +0.04069300f, +0.08257346f, +0.14548601f,
+    -0.18673635f, -0.13438253f, -0.088401966f, -0.05205436f, -0.02032501f, +0.012399545f, +0.05127183f, +0.10316186f,
+    -0.10807011f, -0.065903045f, -0.032206114f, -0.0062006037f, +0.020679146f, +0.04422085f, +0.08313074f, +0.16821936f,
+    -0.22979105f, -0.14431947f, -0.07689272f, -0.02755307f, +0.009225173f, +0.046684854f, +0.08834142f, +0.13766693f,
+    -0.22114082f, -0.12612148f, -0.06890522f, -0.016128855f, +0.03691900f, +0.08474852f, +0.14940020f, +0.23229980f,
+    -0.14933491f, -0.099693604f, -0.06738499f, -0.037100967f, -0.009332986f, +0.023535024f, +0.060272533f, +0.109464675f,
+    -0.20200425f, -0.07398328f, -0.038700905f, -0.01714807f, +0.011161969f, +0.04528101f, +0.08902637f, +0.19573534f,
+    -0.16645233f, -0.124482535f, -0.089342155f, -0.04427387f, -0.007353691f, +0.028033108f, +0.066108435f, +0.15552913f,
+    -0.22295763f, -0.059887577f, -0.018804537f, +0.020141022f, +0.059682943f, +0.097920544f, +0.14080113f, +0.25698325f,
+    -0.14248224f, -0.089685425f, -0.050101686f, -0.017257255f, +0.011412255f, +0.040830314f, +0.07400172f, +0.11997315f,
+    -0.18649384f, -0.113997504f, -0.067775466f, -0.033394672f, +0.006586988f, +0.05312057f, +0.10433043f, +0.22344802f,
+    -0.16138338f, -0.108194515f, -0.07600300f, -0.05135381f, -0.023365447f, +0.0087320795f, +0.045431953f, +0.09113002f,
+    -0.12630440f, -0.07225349f, -0.032280035f, +0.0029231994f, +0.019239848f, +0.05081419f, +0.077840395f, +0.121695265f,
+    -0.08928155f, -0.044983763f, -0.009889568f, +0.020831043f, +0.05684458f, +0.09409702f, +0.13867535f, +0.19084482f,
+    -0.14182915f, -0.11380146f, -0.06904074f, -0.002002765f, +0.034864165f, +0.070399575f, +0.11403063f, +0.15394832f,
+    -0.10876417f, -0.056122433f, -0.02267638f, +0.011113975f, +0.039639056f, +0.074084364f, +0.10155376f, +0.12540291f,
+    -0.17693359f, -0.13940524f, -0.10049578f, -0.06796275f, -0.036915872f, +0.00062823476f, +0.042142134f, +0.17906062f,
+    -0.09253492f, -0.04290128f, -0.006311852f, +0.023908244f, +0.049849935f, +0.078770354f, +0.10818172f, +0.15166481f,
+    -0.12429565f, -0.07392063f, -0.029114135f, +0.0059440783f, +0.042675965f, +0.08425635f, +0.13836108f, +0.18634140f,
+    -0.11795639f, -0.07033707f, -0.034163877f, -0.0008773357f, +0.03334606f, +0.07188203f, +0.12216825f, +0.17097956f,
+    -0.18718453f, -0.14090346f, -0.097799584f, -0.059522875f, -0.019208657f, +0.03079176f, +0.09334672f, +0.15811224f,
+    -0.27198875f, -0.16546582f, -0.11433405f, -0.06933013f, -0.04026183f, -0.0061146915f, +0.029263576f, +0.07322499f,
+    -0.18471734f, -0.102074504f, -0.06492570f, -0.034418534f, -0.009636157f, +0.023043344f, +0.05751496f, +0.09905984f,
+    -0.22826399f, -0.15946552f, -0.09913176f, -0.06585259f, -0.03252090f, +0.001313243f, +0.03556729f, +0.21612854f,
+    -0.13243781f, -0.087299444f, -0.049820945f, -0.016216082f, +0.01799807f, +0.057916876f, +0.09001349f, +0.13221787f,
+    -0.19516511f, -0.120894566f, -0.076130204f, -0.051442243f, -0.029535033f, -0.0020043184f, +0.029452588f, +0.075566076f,
+    -0.27272871f, -0.15841717f, -0.105432935f, -0.06792948f, -0.024532158f, +0.014960791f, +0.054415092f, +0.101517834f,
+    -0.21153601f, -0.15015371f, -0.08676790f, -0.04414934f, -0.0042129597f, +0.033762872f, +0.07589151f, +0.12768789f,
+    -0.090428725f, -0.037582967f, +0.0013173596f, +0.03900247f, +0.06840049f, +0.116906695f, +0.16584939f, +0.25382105f,
+    -0.13446195f, -0.07865091f, -0.039625354f, -0.0028398742f, +0.03019514f, +0.06799379f, +0.11850997f, +0.17521496f,
+    -0.11350345f, -0.058599845f, -0.017512511f, +0.019431496f, +0.055897832f, +0.093173414f, +0.14820710f, +0.22092152f,
+    -0.15165758f, -0.08869354f, -0.04974287f, -0.01705474f, +0.013134752f, +0.04367713f, +0.07733791f, +0.12430801f,
+    -0.09329869f, -0.04673005f, -0.00045857552f, +0.042781368f, +0.07802363f, +0.11887439f, +0.16250038f, +0.28612965f,
+    -0.12571070f, -0.07786012f, -0.03843933f, -0.0075433915f, +0.025822964f, +0.066053316f, +0.12021536f, +0.18341768f,
+    -0.16079275f, -0.04921760f, -0.006114644f, +0.026215268f, +0.05699377f, +0.09813471f, +0.16080129f, +0.23786584f,
+    -0.09980837f, -0.048535258f, -0.0096120685f, +0.025387142f, +0.05979822f, +0.09875251f, +0.14474337f, +0.20324114f,
+    -0.15846540f, -0.09938028f, -0.061492465f, -0.03523542f, -0.0061364113f, +0.024916094f, +0.06037314f, +0.106796466f,
+    -0.20557843f, -0.123237535f, -0.07734871f, -0.044549115f, -0.017114898f, +0.01616654f, +0.049574375f, +0.092319444f,
+    -0.19221115f, -0.14642999f, -0.091701314f, -0.055265956f, -0.021026207f, +0.017720066f, +0.05786183f, +0.110154524f,
+    -0.09956386f, -0.03870283f, +0.003052007f, +0.034851722f, +0.06256365f, +0.09628840f, +0.13979156f, +0.16582295f,
+    -0.18026546f, -0.12448310f, -0.07424377f, -0.03954519f, -0.01221123f, +0.028641058f, +0.100819774f, +0.18240699f,
+    -0.21520759f, -0.15573645f, -0.09820838f, -0.051450998f, -0.012993679f, +0.021135861f, +0.058727216f, +0.105848536f,
+    -0.11207385f, -0.08335689f, -0.048542723f, -0.023198519f, +0.0039304253f, +0.037778318f, +0.07813917f, +0.13106476f,
+    -0.17849164f, -0.120988995f, -0.078016765f, -0.043093704f, -0.016565649f, +0.015182641f, +0.050754096f, +0.09595712f,
+    -0.22132620f, -0.13407415f, -0.065785654f, -0.013291034f, +0.032098345f, +0.07478225f, +0.12431934f, +0.19174045f,
+    -0.095454164f, -0.051898945f, -0.015116375f, -0.012596778f, +0.018636847f, +0.05006925f, +0.087654814f, +0.13754296f,
+    -0.15254061f, -0.09576059f, -0.052086458f, -0.01596074f, +0.017607626f, +0.04778498f, +0.08950204f, +0.14901252f,
+    -0.26057002f, -0.12472382f, -0.074396215f, -0.03764066f, +0.0011168446f, +0.061569117f, +0.10793752f, +0.19771695f,
+    -0.08661132f, -0.045195263f, -0.016098704f, +0.012780116f, +0.040476497f, +0.074102715f, +0.074102715f, +0.12635531f,
+    -0.14047913f, -0.059587404f, -0.016261123f, +0.019801628f, +0.053541403f, +0.096650146f, +0.15005490f, +0.21051759f,
+    -0.22986396f, -0.11964334f, -0.07266585f, -0.026522418f, +0.018169926f, +0.058630653f, +0.100647695f, +0.15919648f,
+    -0.13251697f, -0.077567816f, -0.042766172f, -0.011389967f, +0.01831755f, +0.05304656f, +0.09620367f, +0.15567583f,
+    -0.119819686f, -0.06772876f, -0.028123451f, +0.00876240f, +0.014405836f, +0.048829112f, +0.08422175f, +0.13823749f,
+    -0.16379014f, -0.08956941f, -0.041652776f, +0.008921398f, +0.05473602f, +0.10037984f, +0.16022855f, +0.23457925f,
+    -0.115844205f, -0.05939626f, -0.020390417f, +0.01374377f, +0.044976473f, +0.07873563f, +0.12207942f, +0.18412720f,
+    -0.19048831f, -0.07587487f, -0.03220580f, -0.00011795067f, +0.02721784f, +0.04380719f, +0.07886723f, +0.13193911f,
+    -0.13935551f, -0.092902906f, -0.052706074f, -0.017797327f, +0.015312965f, +0.056098964f, +0.11203423f, +0.24448302f,
+    -0.17986591f, -0.10738580f, -0.06376371f, -0.026595421f, +0.00842492f, +0.04272362f, +0.08608052f, +0.15240218f,
+    -0.10953678f, -0.057022586f, -0.012483291f, +0.024463262f, +0.06076792f, +0.09776234f, +0.12983681f, +0.18648379f,
+    -0.16471463f, -0.089491285f, -0.037574016f, +0.004444791f, +0.039293647f, +0.07845859f, +0.12893885f, +0.23508036f
+};
+
+// --- 2-bit TCQ decode codebook (256-state), verbatim from buun fattn-common.cuh:129-162 ---
+static __constant__ float d_turbo2_tcq_codebook_fattn[256] = {
+    -0.18030643f, -0.11009848f, -0.04742626f, +0.02894132f, -0.10523465f, -0.031312924f, +0.031491395f, +0.12263535f,
+    -0.15660362f, -0.055477407f, +0.0046675834f, +0.06166081f, -0.07506216f, -0.016963918f, +0.043737844f, +0.116496615f,
+    -0.08632783f, -0.022493735f, +0.041032985f, +0.10660284f, -0.06274858f, -0.0036939639f, +0.02095157f, +0.07539709f,
+    -0.09802641f, -0.008419088f, +0.059072323f, +0.17311879f, -0.093109086f, -0.02654333f, +0.014827672f, +0.07793592f,
+    -0.031235758f, +0.01271591f, +0.08752262f, +0.17246453f, -0.14595252f, -0.07227624f, +0.013628688f, +0.08131674f,
+    -0.036909282f, +0.0018896917f, +0.05209119f, +0.12407892f, -0.13689458f, -0.06054520f, +0.0064648795f, +0.07551241f,
+    -0.18980840f, -0.110128626f, -0.046503957f, +0.026387159f, -0.034967307f, +0.04810357f, +0.072072044f, +0.14355458f,
+    -0.10182410f, -0.02907887f, +0.014033012f, +0.083419636f, -0.056140676f, +0.008405868f, +0.066070884f, +0.14037225f,
+    -0.117427245f, -0.047159385f, +0.016928354f, +0.08142885f, -0.029359628f, +0.045608785f, +0.10559447f, +0.20061271f,
+    -0.040425077f, +0.029068163f, +0.08408973f, +0.13628258f, -0.16633821f, -0.10711727f, -0.04196669f, +0.027895834f,
+    -0.0054065837f, +0.058898676f, +0.12688550f, +0.18268861f, -0.16287325f, -0.11218357f, -0.07165227f, -0.009524379f,
+    -0.24026902f, -0.073219374f, -0.0005165726f, +0.05959821f, -0.05532953f, +0.027044486f, +0.09425678f, +0.15356481f,
+    -0.14381111f, -0.10563502f, -0.037867088f, +0.023611993f, -0.03624307f, +0.049588434f, +0.12192037f, +0.23462485f,
+    -0.14990251f, -0.09659304f, -0.05886742f, +0.014878461f, -0.009889551f, +0.06910514f, +0.12120181f, +0.22596690f,
+    -0.08290075f, -0.009009629f, +0.066151775f, +0.12188313f, -0.11591514f, -0.06952189f, -0.031633306f, +0.023740824f,
+    -0.20510401f, -0.103369795f, +0.09148037f, +0.17268716f, -0.16597997f, -0.09207068f, -0.032810967f, +0.024847647f,
+    -0.02487482f, +0.049298953f, +0.09624215f, +0.14217524f, -0.18418685f, -0.10147012f, -0.05841265f, +0.008057022f,
+    -0.14269894f, -0.092456274f, -0.026881337f, +0.049792137f, -0.019881032f, +0.030333601f, +0.09736802f, +0.17764080f,
+    -0.19579841f, -0.114739306f, -0.026823774f, +0.07466014f, -0.09001050f, -0.041468445f, +0.028473806f, +0.08870695f,
+    -0.019396419f, +0.042828932f, +0.10885327f, +0.13335012f, -0.15005013f, -0.074581385f, -0.028608415f, +0.03848942f,
+    -0.09687270f, -0.057059396f, +0.0077843578f, +0.06302297f, -0.23247094f, -0.14509225f, -0.032651436f, +0.027010715f,
+    -0.047595482f, +0.06280303f, +0.114691675f, +0.17124057f, -0.21092793f, -0.13704823f, -0.07340412f, +0.0039013291f,
+    -0.062834196f, +0.012601906f, +0.012601906f, +0.08721347f, -0.13256435f, -0.024173854f, +0.07723171f, +0.14801070f,
+    -0.06471605f, -0.0017903054f, -0.0017903054f, +0.058302354f, -0.09731802f, -0.03400696f, +0.02762442f, +0.08986137f,
+    -0.08288722f, -0.019051429f, +0.045709886f, +0.15211061f, -0.09507891f, -0.015612489f, +0.025347246f, +0.087257534f,
+    -0.066236064f, -0.0047936034f, +0.06386274f, +0.15401669f, -0.105809286f, -0.051802177f, +0.01073050f, +0.08292137f,
+    -0.11884470f, -0.04404144f, +0.02550729f, +0.02550729f, -0.01731189f, +0.062161792f, +0.12127554f, +0.21981733f,
+    -0.17066145f, -0.11660990f, -0.049425896f, +0.021293938f, -0.04711412f, +0.026577346f, +0.055197213f, +0.12541275f,
+    -0.028268812f, +0.015206398f, +0.09002519f, +0.12699963f, -0.10059831f, -0.026676945f, +0.059903253f, +0.13054545f,
+    -0.09582803f, -0.033371232f, +0.010346129f, +0.066766635f, -0.09964944f, -0.028686784f, +0.021184925f, +0.09120017f,
+    -0.16957201f, -0.07594450f, +0.04172865f, +0.18313301f, -0.051526368f, +0.011877304f, +0.011877304f, +0.07956263f,
+    -0.13432936f, -0.05269006f, +0.03536416f, +0.117640756f, -0.022776067f, +0.042032316f, +0.10472976f, +0.18042557f
+};
+
+// TCQ decode-time alphas: per-compilation-unit copies, set by cudaMemcpyToSymbol before kernel launch
+static __constant__ float d_tcq_decode_alpha_k_fattn = 1.0f;
+static __constant__ float d_tcq_decode_alpha_v_fattn = 1.0f;
+
+// TCQ 3-bit K dot product: 9-bit state → codebook lookup
+// Core implementation takes explicit codebook pointer for SMEM/constant flexibility
+template<int D, int nthreads>
+static __device__ __forceinline__ float vec_dot_fattn_vec_KQ_turbo3_tcq_cb(
+    const char * __restrict__ K_c, const void * __restrict__ Q_v,
+    const int * __restrict__ Q_q8, const void * __restrict__ Q_ds_v,
+    const float * __restrict__ cb) {
+    const block_turbo3_tcq * K_tcq = (const block_turbo3_tcq *) K_c;
+    GGML_UNUSED(Q_q8); GGML_UNUSED(Q_ds_v);
+    constexpr int cpy_nb = ggml_cuda_get_max_cpy_bytes();
+    constexpr int cpy_ne = cpy_nb / 4;
+    float sum = 0.0f;
+    int prev_ib = -1;
+    float norm = 0.0f;
+#pragma unroll
+    for (int k_KQ_0 = 0; k_KQ_0 < D/2; k_KQ_0 += nthreads*cpy_ne) {
+        const int base_f2 = k_KQ_0 + (threadIdx.x % nthreads) * cpy_ne;
+        const int elem0 = base_f2 * 2;
+        const int ib = elem0 / QK_TURBO3_TCQ;
+        const int j_start = elem0 % QK_TURBO3_TCQ;
+
+        if (ib != prev_ib) {
+            norm = __half2float(K_tcq[ib].norm);
+            prev_ib = ib;
+        }
+
+#pragma unroll
+        for (int k_KQ_1 = 0; k_KQ_1 < cpy_ne; ++k_KQ_1) {
+            const int lj = k_KQ_1 * 2;
+            const int t0 = j_start + lj;
+            const int t1 = t0 + 1;
+            const int bp0 = t0 * 3;
+            const uint16_t raw0 = (uint16_t)K_tcq[ib].qs[bp0/8] | ((uint16_t)K_tcq[ib].qs[bp0/8 + 1] << 8);
+            const float k0 = cb[(raw0 >> (bp0 % 8)) & 0x1FF] * norm;
+            const int bp1 = t1 * 3;
+            const uint16_t raw1 = (uint16_t)K_tcq[ib].qs[bp1/8] | ((uint16_t)K_tcq[ib].qs[bp1/8 + 1] << 8);
+            const float k1 = cb[(raw1 >> (bp1 % 8)) & 0x1FF] * norm;
+#ifdef V_DOT2_F32_F16_AVAILABLE
+            const float2 qf = __half22float2(((const half2 *) Q_v)[k_KQ_0/nthreads + k_KQ_1]);
+#else
+            const float2 qf = ((const float2 *) Q_v)[k_KQ_0/nthreads + k_KQ_1];
+#endif
+            sum += k0 * qf.x + k1 * qf.y;
+        }
+    }
+    return sum;
+}
+
+// Wrapper using __constant__ codebook (for function pointer dispatch)
+template<int D, int nthreads>
+static __device__ __forceinline__ float vec_dot_fattn_vec_KQ_turbo3_tcq(
+    const char * __restrict__ K_c, const void * __restrict__ Q_v,
+    const int * __restrict__ Q_q8, const void * __restrict__ Q_ds_v) {
+    return vec_dot_fattn_vec_KQ_turbo3_tcq_cb<D, nthreads>(K_c, Q_v, Q_q8, Q_ds_v, d_turbo3_tcq_codebook_fattn);
+}
+
+// TCQ 2-bit K dot product: 8-bit state → codebook lookup
+template<int D, int nthreads>
+static __device__ __forceinline__ float vec_dot_fattn_vec_KQ_turbo2_tcq_cb(
+    const char * __restrict__ K_c, const void * __restrict__ Q_v,
+    const int * __restrict__ Q_q8, const void * __restrict__ Q_ds_v,
+    const float * __restrict__ cb) {
+    const block_turbo2_tcq * K_tcq = (const block_turbo2_tcq *) K_c;
+    GGML_UNUSED(Q_q8); GGML_UNUSED(Q_ds_v);
+    constexpr int cpy_nb = ggml_cuda_get_max_cpy_bytes();
+    constexpr int cpy_ne = cpy_nb / 4;
+    float sum = 0.0f;
+    int prev_ib = -1;
+    float norm = 0.0f;
+#pragma unroll
+    for (int k_KQ_0 = 0; k_KQ_0 < D/2; k_KQ_0 += nthreads*cpy_ne) {
+        const int base_f2 = k_KQ_0 + (threadIdx.x % nthreads) * cpy_ne;
+        const int elem0 = base_f2 * 2;
+        const int ib = elem0 / QK_TURBO2_TCQ;
+        const int j_start = elem0 % QK_TURBO2_TCQ;
+
+        if (ib != prev_ib) {
+            norm = __half2float(K_tcq[ib].norm);
+            prev_ib = ib;
+        }
+
+#pragma unroll
+        for (int k_KQ_1 = 0; k_KQ_1 < cpy_ne; ++k_KQ_1) {
+            const int lj = k_KQ_1 * 2;
+            const int t0 = j_start + lj;
+            const int t1 = t0 + 1;
+            const int bp0 = t0 * 2;
+            const uint16_t raw0 = (uint16_t)K_tcq[ib].qs[bp0/8] | ((uint16_t)K_tcq[ib].qs[bp0/8 + 1] << 8);
+            const float k0 = cb[(raw0 >> (bp0 % 8)) & 0xFF] * norm;
+            const int bp1 = t1 * 2;
+            const uint16_t raw1 = (uint16_t)K_tcq[ib].qs[bp1/8] | ((uint16_t)K_tcq[ib].qs[bp1/8 + 1] << 8);
+            const float k1 = cb[(raw1 >> (bp1 % 8)) & 0xFF] * norm;
+#ifdef V_DOT2_F32_F16_AVAILABLE
+            const float2 qf = __half22float2(((const half2 *) Q_v)[k_KQ_0/nthreads + k_KQ_1]);
+#else
+            const float2 qf = ((const float2 *) Q_v)[k_KQ_0/nthreads + k_KQ_1];
+#endif
+            sum += k0 * qf.x + k1 * qf.y;
+        }
+    }
+    return sum;
+}
+
+// Wrapper using __constant__ codebook (for function pointer dispatch)
+template<int D, int nthreads>
+static __device__ __forceinline__ float vec_dot_fattn_vec_KQ_turbo2_tcq(
+    const char * __restrict__ K_c, const void * __restrict__ Q_v,
+    const int * __restrict__ Q_q8, const void * __restrict__ Q_ds_v) {
+    return vec_dot_fattn_vec_KQ_turbo2_tcq_cb<D, nthreads>(K_c, Q_v, Q_q8, Q_ds_v, d_turbo2_tcq_codebook_fattn);
+}
+
+// TCQ 3-bit V dequant: 9-bit state → codebook lookup
+// Core implementation takes explicit codebook pointer for SMEM/constant flexibility
+template <typename T, int ne>
+static __device__ __forceinline__ void dequantize_V_turbo3_tcq_cb(
+        const void * __restrict__ vx, void * __restrict__ dst, const int64_t i0,
+        const float * __restrict__ cb) {
+    const block_turbo3_tcq * x = (const block_turbo3_tcq *) vx;
+    const int64_t ib = i0 / QK_TURBO3_TCQ;
+    const int     j0 = (int)(i0 % QK_TURBO3_TCQ);
+    const float norm = __half2float(x[ib].norm) * d_tcq_decode_alpha_v_fattn;
+    static_assert(ne == 2 || ne == 4 || ne == 8, "bad ne");
+    float vals[ne];
+#pragma unroll
+    for (int l = 0; l < ne; l++) {
+        const int t = j0 + l;
+        const int bit_pos = t * 3;
+        const uint16_t raw = (uint16_t)x[ib].qs[bit_pos/8] | ((uint16_t)x[ib].qs[bit_pos/8 + 1] << 8);
+        const int state = (raw >> (bit_pos % 8)) & 0x1FF;
+        vals[l] = cb[state] * norm;
+    }
+#ifdef FP16_AVAILABLE
+    if constexpr (std::is_same_v<T, half>) {
+        for (int l0 = 0; l0 < ne; l0 += 2)
+            ((half2 *)dst)[l0/2] = make_half2(__float2half(vals[l0]), __float2half(vals[l0+1]));
+    } else
+#endif
+    if constexpr (std::is_same_v<T, float>) {
+        for (int l = 0; l < ne; ++l) ((float *)dst)[l] = vals[l];
+    } else { static_assert(std::is_same_v<T, void>, "bad type"); }
+}
+
+// Wrapper using __constant__ codebook (for function pointer dispatch via dequantize_V_t)
+template <typename T, int ne>
+static __device__ __forceinline__ void dequantize_V_turbo3_tcq(
+        const void * __restrict__ vx, void * __restrict__ dst, const int64_t i0) {
+    dequantize_V_turbo3_tcq_cb<T, ne>(vx, dst, i0, d_turbo3_tcq_codebook_fattn);
+}
+
+// TCQ 2-bit V dequant: 8-bit state → codebook lookup
+template <typename T, int ne>
+static __device__ __forceinline__ void dequantize_V_turbo2_tcq_cb(
+        const void * __restrict__ vx, void * __restrict__ dst, const int64_t i0,
+        const float * __restrict__ cb) {
+    const block_turbo2_tcq * x = (const block_turbo2_tcq *) vx;
+    const int64_t ib = i0 / QK_TURBO2_TCQ;
+    const int     j0 = (int)(i0 % QK_TURBO2_TCQ);
+    const float norm = __half2float(x[ib].norm) * d_tcq_decode_alpha_v_fattn;
+    static_assert(ne == 2 || ne == 4 || ne == 8, "bad ne");
+    float vals[ne];
+#pragma unroll
+    for (int l = 0; l < ne; l++) {
+        const int t = j0 + l;
+        const int bit_pos = t * 2;
+        const uint16_t raw = (uint16_t)x[ib].qs[bit_pos/8] | ((uint16_t)x[ib].qs[bit_pos/8 + 1] << 8);
+        const int state = (raw >> (bit_pos % 8)) & 0xFF;
+        vals[l] = cb[state] * norm;
+    }
+#ifdef FP16_AVAILABLE
+    if constexpr (std::is_same_v<T, half>) {
+        for (int l0 = 0; l0 < ne; l0 += 2)
+            ((half2 *)dst)[l0/2] = make_half2(__float2half(vals[l0]), __float2half(vals[l0+1]));
+    } else
+#endif
+    if constexpr (std::is_same_v<T, float>) {
+        for (int l = 0; l < ne; ++l) ((float *)dst)[l] = vals[l];
+    } else { static_assert(std::is_same_v<T, void>, "bad type"); }
+}
+
+// Wrapper using __constant__ codebook (for function pointer dispatch via dequantize_V_t)
+template <typename T, int ne>
+static __device__ __forceinline__ void dequantize_V_turbo2_tcq(
+        const void * __restrict__ vx, void * __restrict__ dst, const int64_t i0) {
+    dequantize_V_turbo2_tcq_cb<T, ne>(vx, dst, i0, d_turbo2_tcq_codebook_fattn);
+}
+
diff --git a/llama.cpp/ggml/src/ggml-cuda/fattn-vec.cuh b/llama.cpp/ggml/src/ggml-cuda/fattn-vec.cuh
index b72c00bd9..abcce887c 100644
--- a/llama.cpp/ggml/src/ggml-cuda/fattn-vec.cuh
+++ b/llama.cpp/ggml/src/ggml-cuda/fattn-vec.cuh
@@ -79,7 +79,11 @@ static __global__ void flash_attn_ext_vec(
     // opencoti-hook: turboquant perf-lift (Level B) — F5 M6-S2 #392. Turbo K reads centroid×norm
     // against a FLOAT Q (Q_q8_1=false below) and its vec_dot mirrors the f16 path, so it takes the
     // f16-style KQ thread split. Turbo V uses the quantized split (nthreads_V_q, V_rows_per_thread=4).
-    constexpr bool is_turbo_K = (type_K == GGML_TYPE_TURBO2_0 || type_K == GGML_TYPE_TURBO3_0 || type_K == GGML_TYPE_TURBO4_0);
+    constexpr bool is_turbo_K = (type_K == GGML_TYPE_TURBO2_0 || type_K == GGML_TYPE_TURBO3_0 || type_K == GGML_TYPE_TURBO4_0
+        // opencoti-hook: TCQ FA-VEC decode (#447/#500). TCQ K is FWHT-rotated like turbo_0 — it must take
+        // the f16-style KQ split AND keep Q FLOAT (Q_q8_1=false below); q8_1-quantizing Q destroys the
+        // WHT-rotation precision the trellis reconstruction needs. V uses our quantized split (ne=4), same as turbo_0.
+        || type_K == GGML_TYPE_TURBO3_TCQ || type_K == GGML_TYPE_TURBO2_TCQ);
     constexpr int nthreads    = ggml_cuda_fattn_vec_get_nthreads_device();
     constexpr int nthreads_KQ = (type_K == GGML_TYPE_F16 || type_K == GGML_TYPE_BF16 || is_turbo_K) ? 128 / cpy_nb : nthreads_KQ_q;
     constexpr int nthreads_V  = (type_V == GGML_TYPE_F16 || type_V == GGML_TYPE_BF16) ? 128 / cpy_nb : nthreads_V_q;
@@ -588,6 +592,7 @@ void ggml_cuda_flash_attn_ext_vec_case(ggml_backend_cuda_context & ctx, ggml_ten
     extern DECL_FATTN_VEC_CASE(D, type_K, GGML_TYPE_Q4_0); \
     extern DECL_FATTN_VEC_CASE(D, type_K, GGML_TYPE_Q4_1); \
     extern DECL_FATTN_VEC_CASE(D, type_K, GGML_TYPE_Q5_0); \
+    extern DECL_FATTN_VEC_CASE(D, type_K, GGML_TYPE_Q6_0); /* opencoti-hook: q6_0 KV (Anbeeld port) */ \
     extern DECL_FATTN_VEC_CASE(D, type_K, GGML_TYPE_Q5_1); \
     extern DECL_FATTN_VEC_CASE(D, type_K, GGML_TYPE_Q8_0); \
     extern DECL_FATTN_VEC_CASE(D, type_K, GGML_TYPE_BF16); \
@@ -596,6 +601,7 @@ EXTERN_DECL_FATTN_VEC_CASES( 64, GGML_TYPE_F16)
 EXTERN_DECL_FATTN_VEC_CASES( 64, GGML_TYPE_Q4_0)
 EXTERN_DECL_FATTN_VEC_CASES( 64, GGML_TYPE_Q4_1)
 EXTERN_DECL_FATTN_VEC_CASES( 64, GGML_TYPE_Q5_0)
+EXTERN_DECL_FATTN_VEC_CASES( 64, GGML_TYPE_Q6_0) // opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port)
 EXTERN_DECL_FATTN_VEC_CASES( 64, GGML_TYPE_Q5_1)
 EXTERN_DECL_FATTN_VEC_CASES( 64, GGML_TYPE_Q8_0)
 EXTERN_DECL_FATTN_VEC_CASES( 64, GGML_TYPE_BF16)
@@ -604,6 +610,7 @@ EXTERN_DECL_FATTN_VEC_CASES(128, GGML_TYPE_F16)
 EXTERN_DECL_FATTN_VEC_CASES(128, GGML_TYPE_Q4_0)
 EXTERN_DECL_FATTN_VEC_CASES(128, GGML_TYPE_Q4_1)
 EXTERN_DECL_FATTN_VEC_CASES(128, GGML_TYPE_Q5_0)
+EXTERN_DECL_FATTN_VEC_CASES(128, GGML_TYPE_Q6_0) // opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port)
 EXTERN_DECL_FATTN_VEC_CASES(128, GGML_TYPE_Q5_1)
 EXTERN_DECL_FATTN_VEC_CASES(128, GGML_TYPE_Q8_0)
 EXTERN_DECL_FATTN_VEC_CASES(128, GGML_TYPE_BF16)
@@ -612,6 +619,7 @@ EXTERN_DECL_FATTN_VEC_CASES(256, GGML_TYPE_F16)
 EXTERN_DECL_FATTN_VEC_CASES(256, GGML_TYPE_Q4_0)
 EXTERN_DECL_FATTN_VEC_CASES(256, GGML_TYPE_Q4_1)
 EXTERN_DECL_FATTN_VEC_CASES(256, GGML_TYPE_Q5_0)
+EXTERN_DECL_FATTN_VEC_CASES(256, GGML_TYPE_Q6_0) // opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port)
 EXTERN_DECL_FATTN_VEC_CASES(256, GGML_TYPE_Q5_1)
 EXTERN_DECL_FATTN_VEC_CASES(256, GGML_TYPE_Q8_0)
 EXTERN_DECL_FATTN_VEC_CASES(256, GGML_TYPE_BF16)
@@ -626,11 +634,27 @@ EXTERN_DECL_FATTN_VEC_CASES(256, GGML_TYPE_BF16)
 EXTERN_DECL_FATTN_VEC_TURBO(GGML_TYPE_TURBO2_0)
 EXTERN_DECL_FATTN_VEC_TURBO(GGML_TYPE_TURBO3_0)
 EXTERN_DECL_FATTN_VEC_TURBO(GGML_TYPE_TURBO4_0)
+// opencoti-hook: TCQ FA-VEC decode (#447/#500) — trellis-coded turbo K==V vec instances (64/128/256).
+EXTERN_DECL_FATTN_VEC_TURBO(GGML_TYPE_TURBO3_TCQ)
+EXTERN_DECL_FATTN_VEC_TURBO(GGML_TYPE_TURBO2_TCQ)
+// opencoti-hook: asymmetric TCQ (#513) — turbo3_tcq/turbo2_tcq + reverse, head_dim {64,128,256}.
+extern DECL_FATTN_VEC_CASE( 64, GGML_TYPE_TURBO3_TCQ, GGML_TYPE_TURBO2_TCQ);
+extern DECL_FATTN_VEC_CASE(128, GGML_TYPE_TURBO3_TCQ, GGML_TYPE_TURBO2_TCQ);
+extern DECL_FATTN_VEC_CASE(256, GGML_TYPE_TURBO3_TCQ, GGML_TYPE_TURBO2_TCQ);
+extern DECL_FATTN_VEC_CASE( 64, GGML_TYPE_TURBO2_TCQ, GGML_TYPE_TURBO3_TCQ);
+extern DECL_FATTN_VEC_CASE(128, GGML_TYPE_TURBO2_TCQ, GGML_TYPE_TURBO3_TCQ);
+extern DECL_FATTN_VEC_CASE(256, GGML_TYPE_TURBO2_TCQ, GGML_TYPE_TURBO3_TCQ);
 
 // opencoti #411 (F5 M6-S2 Level-B): head_dim 512 fused VEC for turbo2/3 ONLY — Gemma-4 GLOBAL
 // layers (key_length=512). Definitions in the turbo2/turbo3 instance .cu files. turbo4 stays ≤256.
 extern DECL_FATTN_VEC_CASE(512, GGML_TYPE_TURBO2_0, GGML_TYPE_TURBO2_0);
 extern DECL_FATTN_VEC_CASE(512, GGML_TYPE_TURBO3_0, GGML_TYPE_TURBO3_0);
+// opencoti-hook: TCQ FA-VEC decode (#447/#500) — head_dim 512 trellis VEC for Gemma-4 GLOBAL layers.
+extern DECL_FATTN_VEC_CASE(512, GGML_TYPE_TURBO3_TCQ, GGML_TYPE_TURBO3_TCQ);
+extern DECL_FATTN_VEC_CASE(512, GGML_TYPE_TURBO2_TCQ, GGML_TYPE_TURBO2_TCQ);
+// opencoti-hook: asymmetric TCQ (#513) — head_dim 512 (Gemma-4 GLOBAL), strong-K + reverse.
+extern DECL_FATTN_VEC_CASE(512, GGML_TYPE_TURBO3_TCQ, GGML_TYPE_TURBO2_TCQ);
+extern DECL_FATTN_VEC_CASE(512, GGML_TYPE_TURBO2_TCQ, GGML_TYPE_TURBO3_TCQ);
 
 // opencoti #493 (F5 M6-S4): head_dim 512 f16 VEC — Gemma-4 GLOBAL layer at gqa_ratio<=2 (the E4B
 // assistant-MTP drafter), where MMA/TILE abort (DKQ>256 && gqa<=2, bug-567). The VEC kernel tiles
diff --git a/llama.cpp/ggml/src/ggml-cuda/fattn.cu b/llama.cpp/ggml/src/ggml-cuda/fattn.cu
index c0a964d29..a48785e89 100644
--- a/llama.cpp/ggml/src/ggml-cuda/fattn.cu
+++ b/llama.cpp/ggml/src/ggml-cuda/fattn.cu
@@ -347,6 +347,18 @@ static void ggml_cuda_flash_attn_ext_vec(ggml_backend_cuda_context & ctx, ggml_t
     FATTN_VEC_CASES_ALL_D(GGML_TYPE_BF16, GGML_TYPE_BF16)
 #endif // GGML_CUDA_FA_ALL_QUANTS
 
+    // opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port). The Anbeeld instance set (the 5 K/V combos
+    // ik_llama.cpp ships for q6_0 KV), registered OUTSIDE the FA_ALL_QUANTS block so they compile in
+    // the DEFAULT build — the matching instance .cu files are added unconditionally in
+    // collect_gpu_sources (build-functions.sh section 3c), like the turbo K==V instances below.
+    // q6_0-q6_0 is the symmetric (default-reachable) cell; the 4 mixed cells are reachable only under
+    // FA_ALL_QUANTS (the K!=V early-out in ggml_cuda_get_best_fattn_kernel), same as q5_0's mixed pairs.
+    FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q6_0, GGML_TYPE_Q6_0)
+    FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q8_0, GGML_TYPE_Q6_0)
+    FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q6_0, GGML_TYPE_Q5_0)
+    FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q6_0, GGML_TYPE_F16)
+    FATTN_VEC_CASES_ALL_D(GGML_TYPE_F16,  GGML_TYPE_Q6_0)
+
     // opencoti-hook: turboquant perf-lift (Level B) — F5 M6-S2 #392. Fused turbo K/V vec path
     // (always registered, independent of GGML_CUDA_FA_ALL_QUANTS). turbo2/3/4 at head_dim ∈
     // {64,128,256}; turbo8 / head_dim 512 are NOT here (Level-A materialize lift handles them).
@@ -354,11 +366,26 @@ static void ggml_cuda_flash_attn_ext_vec(ggml_backend_cuda_context & ctx, ggml_t
     FATTN_VEC_CASES_ALL_D(GGML_TYPE_TURBO3_0, GGML_TYPE_TURBO3_0)
     FATTN_VEC_CASES_ALL_D(GGML_TYPE_TURBO4_0, GGML_TYPE_TURBO4_0)
 
+    // opencoti-hook: TCQ FA-VEC decode (#447/#500) — trellis turbo K==V at head_dim {64,128,256}.
+    FATTN_VEC_CASES_ALL_D(GGML_TYPE_TURBO3_TCQ, GGML_TYPE_TURBO3_TCQ)
+    FATTN_VEC_CASES_ALL_D(GGML_TYPE_TURBO2_TCQ, GGML_TYPE_TURBO2_TCQ)
+    // opencoti-hook: asymmetric TCQ (#513) — turbo3_tcq K / turbo2_tcq V (strong-K) + reverse, head_dim {64,128,256}.
+    // K!=V both-TCQ: graph fuses (build_attn_mha) + selector K!=V gate relaxed above. Codebooks baked per-TU
+    // in the instance files (A′ override via tcq{3,2}_set_decode_codebook_{t3t2,t2t3} in set-rows.cu).
+    FATTN_VEC_CASES_ALL_D(GGML_TYPE_TURBO3_TCQ, GGML_TYPE_TURBO2_TCQ)
+    FATTN_VEC_CASES_ALL_D(GGML_TYPE_TURBO2_TCQ, GGML_TYPE_TURBO3_TCQ)
+
     // opencoti #411 (F5 M6-S2 Level-B): head_dim 512 fused VEC for turbo2/3 ONLY (Gemma-4 GLOBAL
     // layers). turbo4/turbo8 keep no D=512 instance (Level-A materialize). Matches the fattn.cu
     // kernel-selection relax + the L3 graph gate (both turbo2/3 + D==512 only).
     FATTN_VEC_CASE(512, GGML_TYPE_TURBO2_0, GGML_TYPE_TURBO2_0)
     FATTN_VEC_CASE(512, GGML_TYPE_TURBO3_0, GGML_TYPE_TURBO3_0)
+    // opencoti-hook: TCQ FA-VEC decode (#447/#500) — head_dim 512 trellis for Gemma-4 GLOBAL layers.
+    FATTN_VEC_CASE(512, GGML_TYPE_TURBO3_TCQ, GGML_TYPE_TURBO3_TCQ)
+    FATTN_VEC_CASE(512, GGML_TYPE_TURBO2_TCQ, GGML_TYPE_TURBO2_TCQ)
+    // opencoti-hook: asymmetric TCQ (#513) — head_dim 512 (Gemma-4 GLOBAL), strong-K + reverse.
+    FATTN_VEC_CASE(512, GGML_TYPE_TURBO3_TCQ, GGML_TYPE_TURBO2_TCQ)
+    FATTN_VEC_CASE(512, GGML_TYPE_TURBO2_TCQ, GGML_TYPE_TURBO3_TCQ)
 
     // opencoti #493 (F5 M6-S4): head_dim 512 f16 VEC dispatch — the runtime counterpart to the
     // selector's D512/gqa<=2 f16 VEC route (ggml_cuda_get_best_fattn_kernel) + the f16-f16 D512
@@ -471,8 +498,14 @@ static best_fattn_kernel ggml_cuda_get_best_fattn_kernel(const int device, const
     }
 
 #ifndef GGML_CUDA_FA_ALL_QUANTS
-    if (K->type != V->type) {
-        return BEST_FATTN_KERNEL_NONE;
+    // opencoti-hook: asymmetric TCQ (#513) — both-TCQ K!=V (turbo3_tcq/turbo2_tcq + reverse) has a
+    // dedicated in-register FA-VEC instance, so let it through; every other K!=V still falls back.
+    {
+        const bool both_tcq = (K->type == GGML_TYPE_TURBO3_TCQ || K->type == GGML_TYPE_TURBO2_TCQ) &&
+                              (V->type == GGML_TYPE_TURBO3_TCQ || V->type == GGML_TYPE_TURBO2_TCQ);
+        if (K->type != V->type && !both_tcq) {
+            return BEST_FATTN_KERNEL_NONE;
+        }
     }
 #endif // GGML_CUDA_FA_ALL_QUANTS
 
@@ -488,11 +521,16 @@ static best_fattn_kernel ggml_cuda_get_best_fattn_kernel(const int device, const
 #endif // GGML_CUDA_FA_ALL_QUANTS
         case GGML_TYPE_Q4_0:
         case GGML_TYPE_Q8_0:
+        case GGML_TYPE_Q6_0: // opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port) — the 5 Anbeeld
+                             // q6_0 vec instances compile in the DEFAULT build (collect_gpu_sources
+                             // section 3c), so q6_0 K is allowed unconditionally like q4_0/q8_0/bf16.
         case GGML_TYPE_BF16:
             break;
         case GGML_TYPE_TURBO2_0: // opencoti-hook: turboquant perf-lift (Level B) — F5 M6-S2 #392
         case GGML_TYPE_TURBO3_0:
         case GGML_TYPE_TURBO4_0:
+        case GGML_TYPE_TURBO3_TCQ: // opencoti-hook: TCQ FA-VEC decode (#447/#500)
+        case GGML_TYPE_TURBO2_TCQ:
             break;
         default:
             return BEST_FATTN_KERNEL_NONE;
@@ -510,13 +548,15 @@ static best_fattn_kernel ggml_cuda_get_best_fattn_kernel(const int device, const
     // VEC kernel (centroid×norm read in-register); never route them to MMA/WMMA/TILE, which can't
     // read turbo blocks. The graph (build_attn_mha, L3) only emits turbo K/V to FLASH_ATTN_EXT when
     // it intends this path (head_dim ≤ 256, decode); turbo8 / head_dim 512 take the materialize lift.
-    if (K->type == GGML_TYPE_TURBO2_0 || K->type == GGML_TYPE_TURBO3_0 || K->type == GGML_TYPE_TURBO4_0) {
+    if (K->type == GGML_TYPE_TURBO2_0 || K->type == GGML_TYPE_TURBO3_0 || K->type == GGML_TYPE_TURBO4_0 ||
+        K->type == GGML_TYPE_TURBO3_TCQ || K->type == GGML_TYPE_TURBO2_TCQ) { // opencoti-hook: TCQ FA-VEC decode (#447/#500)
         // opencoti #411 (F5 M6-S2 Level-B): turbo2/3 also ship a D=512 fused VEC instance for Gemma-4
         // GLOBAL layers (key_length=512) so the 256k global cache reads in-register instead of the
         // Level-A turbo→f16 materialize. turbo4/turbo8 have no D=512 instance → keep the ≤256 cap
         // (Level-A fallback at 512). The L3 graph gate only emits turbo2/3 K/V to FA at 512 to match.
         const bool turbo512_vec = Q->ne[0] == 512 && Q->ne[0] % 64 == 0 && K->ne[1] % FATTN_KQ_STRIDE == 0 &&
-            (K->type == GGML_TYPE_TURBO2_0 || K->type == GGML_TYPE_TURBO3_0);
+            (K->type == GGML_TYPE_TURBO2_0 || K->type == GGML_TYPE_TURBO3_0 ||
+             K->type == GGML_TYPE_TURBO3_TCQ || K->type == GGML_TYPE_TURBO2_TCQ); // opencoti-hook: TCQ #447/#500 (D512 trellis)
         return (can_use_vector_kernel || turbo512_vec) ? BEST_FATTN_KERNEL_VEC : BEST_FATTN_KERNEL_NONE;
     }
 
diff --git a/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu b/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu
index 672bdb5ab..a6ef9eef8 100644
--- a/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu
+++ b/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu
@@ -5391,8 +5391,12 @@ static bool GGML_CALL ggml_backend_cuda_device_supports_op(ggml_backend_dev_t de
                 return (op->type == GGML_TYPE_F32 || op->type == GGML_TYPE_F16 || op->type == GGML_TYPE_BF16 ||
                        op->type == GGML_TYPE_Q4_0 || op->type == GGML_TYPE_Q4_1 || op->type == GGML_TYPE_Q5_0 ||
                        op->type == GGML_TYPE_Q5_1 || op->type == GGML_TYPE_Q8_0 || op->type == GGML_TYPE_IQ4_NL ||
+                       // opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port) — KV-write dst (set-rows.cu generic quant kernel)
+                       op->type == GGML_TYPE_Q6_0 ||
                        op->type == GGML_TYPE_TURBO3_0 || op->type == GGML_TYPE_TURBO4_0 ||
-                       op->type == GGML_TYPE_TURBO2_0 || op->type == GGML_TYPE_TURBO8_0) &&
+                       op->type == GGML_TYPE_TURBO2_0 || op->type == GGML_TYPE_TURBO8_0 ||
+                       // opencoti-hook: TCQ (F5 M6 C5) — turbo3_tcq/turbo2_tcq Viterbi encode (set-rows.cu)
+                       op->type == GGML_TYPE_TURBO3_TCQ || op->type == GGML_TYPE_TURBO2_TCQ) &&
                        op->src[0]->type == GGML_TYPE_F32 &&
                        (op->src[1]->type == GGML_TYPE_I64 || op->src[1]->type == GGML_TYPE_I32);
             } break;
@@ -5445,6 +5449,13 @@ static bool GGML_CALL ggml_backend_cuda_device_supports_op(ggml_backend_dev_t de
                 if (src0_type == GGML_TYPE_Q5_0 && src1_type == GGML_TYPE_F32) {
                     return true;
                 }
+                // opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port) — f32<->q6_0 cpy (mirrors q5_0).
+                if (src0_type == GGML_TYPE_F32 && src1_type == GGML_TYPE_Q6_0) {
+                    return true;
+                }
+                if (src0_type == GGML_TYPE_Q6_0 && src1_type == GGML_TYPE_F32) {
+                    return true;
+                }
                 if (src0_type == GGML_TYPE_F32 && src1_type == GGML_TYPE_Q5_1) {
                     return true;
                 }
diff --git a/llama.cpp/ggml/src/ggml-cuda/set-rows.cu b/llama.cpp/ggml/src/ggml-cuda/set-rows.cu
index 3e4a76076..f0bb7cba2 100644
--- a/llama.cpp/ggml/src/ggml-cuda/set-rows.cu
+++ b/llama.cpp/ggml/src/ggml-cuda/set-rows.cu
@@ -3,6 +3,9 @@
 #include "turbo-dequant.cuh" // opencoti-hook: turboquant — see docs/features/poly_kv.md (M6-S2 Gate-2 write path)
 #include "turbo-innerq.cuh"  // opencoti-hook: turboquant InnerQ (M6-S2) — per-channel equalization
 #include "turbo-innerq-dev.cuh" // opencoti-hook: turboquant perf-lift (M6-S2 Level A) — cross-TU device-scale accessor
+#include "turbo-tcq-cuda.cuh"  // opencoti-hook: TCQ (F5 M6 C5) — Viterbi encode kernels + codebooks + GET_ROWS dequant
+#include "tcq-decode-cb.cuh"   // opencoti-hook: TCQ A′ — cross-TU decode-codebook setters (symmetric override)
+#include <cstring>             // strncmp (cache_k_ vs cache_v_ detection for TCQ alpha)
 #include <cstdlib>
 #include <cmath>
 
@@ -598,6 +601,272 @@ static __global__ void k_set_rows_turbo3(
     if (elem_in_block == 0) blk->norm = __float2half(corrected_norm);
 }
 
+// ===================== TCQ (F5 M6 C5): Trellis-Coded Quantization encode dispatch =====================
+// opencoti-hook: TCQ — drives the buun Viterbi encode kernels (turbo-tcq-cuda.cuh). Ported from buun's
+// ggml_cuda_op_set_rows TURBO*_TCQ branches, adapted to our per-type helper style. The fp32 path is the
+// correctness anchor (byte-faithful to buun); fp64-selectable encode lands as a follow-on.
+
+// Per-device Viterbi backtrace fallback buffer. Only used when the 8 KB/block shared-memory backtrace is
+// unavailable; on cc>=7.5 (e.g. RTX 3090, 99 KB optin) the shared-bt path is the default, so this rarely allocs.
+static uint8_t * tcq_bt_buf   [GGML_CUDA_MAX_DEVICES] = {};
+static size_t    tcq_bt_buf_sz[GGML_CUDA_MAX_DEVICES] = {};
+static void ensure_tcq_bt_buf(int device, size_t bytes) {
+    if (tcq_bt_buf_sz[device] >= bytes) return;
+    if (tcq_bt_buf[device]) CUDA_CHECK(cudaFree(tcq_bt_buf[device]));
+    CUDA_CHECK(cudaMalloc((void**)&tcq_bt_buf[device], bytes));
+    tcq_bt_buf_sz[device] = bytes;
+}
+
+// Context-adaptive norm-scaling override. Sets the turbo-tcq-cuda.cuh __constant__ symbols for THIS TU.
+// Env now (TURBO_TCQ_ALPHA / _ALPHA_V); Stage-4 replaces with the closed-form n_kv schedule.
+static void load_tcq_norm_alpha(int device) {
+    (void) device;
+    const char * sa = getenv("TURBO_TCQ_ALPHA");
+    const char * sv = getenv("TURBO_TCQ_ALPHA_V");
+    if (sa) { float a = (float) atof(sa); CUDA_CHECK(cudaMemcpyToSymbol(d_tcq_norm_alpha,   &a, sizeof(float))); }
+    if (sv) { float a = (float) atof(sv); CUDA_CHECK(cudaMemcpyToSymbol(d_tcq_norm_alpha_v, &a, sizeof(float))); }
+}
+
+// opencoti-hook: TCQ fp64-selectable encode — see docs/evaluations/context.md §5a-TCQ. When TURBO_TCQ_FP64=1,
+// the host dispatches the <idx_t,double> instantiation of the Viterbi encode kernels (acc_t=double), so the
+// path-cost accumulation + norm/whiten reductions run in fp64 — the "true optimal path" oracle for the
+// precision-sensitivity question. Default off ⇒ acc_t=float ⇒ byte-identical to the buun fp32 trellis.
+// Process-global (the env is read once); decode stays fp32 (it reads stored integer states).
+static bool tcq_fp64_enabled() {
+    static int v = -1;
+    if (v < 0) {
+        const char * e = getenv("TURBO_TCQ_FP64");
+        v = (e && atoi(e) != 0) ? 1 : 0;
+        if (v) fprintf(stderr, "TCQ encode: fp64 Viterbi path ENABLED (TURBO_TCQ_FP64=1)\n");
+    }
+    return v != 0;
+}
+
+// opencoti-hook: TCQ activation-dump (F-D, plan lexical-gliding-flurry / #515) — see docs/evaluations/context.md §5a-TCQ.
+// Host side of the dump that k_set_rows_turbo3_tcq's d_tcq_dump_buf write feeds. When TURBO_TCQ_DUMP_ACTS=<file>
+// is set, after each turbo3_tcq encode launch we copy the staged post-FWHT unit-norm 128-vecs D2H and append them
+// as raw little-endian float32 (128 floats/vec) to <file> — caller MUST point this at persistent disk, NEVER /tmp.
+// Capped at TURBO_TCQ_DUMP_MAX vecs (default 200000 ≈ 100 MB). These whitened vecs are codebook-resolution-
+// independent (FWHT precedes the trellis), so this base-tier dump trains the 512-entry turbo3_tcq F-E
+// codebook. Zero overhead when the env is unset: the gate short-circuits and d_tcq_dump_buf stays nullptr (the
+// kernel branch is then a uniform no-op).
+struct tcq_dump_state {
+    bool    checked   = false;
+    bool    active    = false;
+    FILE *  fp        = nullptr;
+    long    max_vecs  = 200000;
+    long    written   = 0;
+    float * dev       = nullptr;   // device staging (grow-only); freed at cap
+    long    dev_vecs  = 0;         // device capacity in vecs
+    float * host      = nullptr;   // pinned host staging (grow-only); freed at cap
+    long    host_vecs = 0;
+};
+static tcq_dump_state g_tcq_dump;
+
+// One-time: read env + open the append file. Idempotent (guarded by `checked`).
+static void tcq_dump_init() {
+    if (g_tcq_dump.checked) return;
+    g_tcq_dump.checked = true;
+    const char * path = getenv("TURBO_TCQ_DUMP_ACTS");
+    if (!path || !path[0]) return;
+    const char * mx = getenv("TURBO_TCQ_DUMP_MAX");
+    if (mx) { long v = atol(mx); if (v > 0) g_tcq_dump.max_vecs = v; }
+    g_tcq_dump.fp = fopen(path, "ab");   // append-binary
+    if (!g_tcq_dump.fp) { fprintf(stderr, "TCQ activation-dump: FAILED to open %s for append\n", path); return; }
+    g_tcq_dump.active = true;
+    fprintf(stderr, "TCQ activation-dump: ENABLED -> %s (cap %ld vecs, 128 f32 each)\n", path, g_tcq_dump.max_vecs);
+}
+
+// Per-launch (BEFORE the kernel): grow the device+host staging to hold `groups` vecs and return the device pointer
+// to bind to d_tcq_dump_buf, or nullptr when the dump is inactive / cap reached (caller then leaves the symbol null).
+static float * tcq_dump_pre_launch(long groups) {
+    if (!g_tcq_dump.active || groups <= 0 || g_tcq_dump.written >= g_tcq_dump.max_vecs) return nullptr;
+    if (g_tcq_dump.dev_vecs < groups) {
+        if (g_tcq_dump.dev) CUDA_CHECK(cudaFree(g_tcq_dump.dev));
+        CUDA_CHECK(cudaMalloc((void**)&g_tcq_dump.dev, (size_t)groups * 128 * sizeof(float)));
+        g_tcq_dump.dev_vecs = groups;
+    }
+    if (g_tcq_dump.host_vecs < groups) {
+        if (g_tcq_dump.host) CUDA_CHECK(cudaFreeHost(g_tcq_dump.host));
+        CUDA_CHECK(cudaMallocHost((void**)&g_tcq_dump.host, (size_t)groups * 128 * sizeof(float)));
+        g_tcq_dump.host_vecs = groups;
+    }
+    return g_tcq_dump.dev;
+}
+
+// Per-launch (AFTER the kernel): copy the staged vecs D2H on `stream`, append (capped) to the file, advance the
+// counter. At the cap, null the symbol + free staging so further launches are zero-overhead no-ops.
+static void tcq_dump_post_launch(long groups, cudaStream_t stream) {
+    if (!g_tcq_dump.active || g_tcq_dump.dev == nullptr || groups <= 0 || g_tcq_dump.written >= g_tcq_dump.max_vecs) return;
+    long take = groups;
+    if (g_tcq_dump.written + take > g_tcq_dump.max_vecs) take = g_tcq_dump.max_vecs - g_tcq_dump.written;
+    CUDA_CHECK(cudaMemcpyAsync(g_tcq_dump.host, g_tcq_dump.dev, (size_t)take * 128 * sizeof(float),
+                               cudaMemcpyDeviceToHost, stream));
+    CUDA_CHECK(cudaStreamSynchronize(stream));
+    fwrite(g_tcq_dump.host, sizeof(float), (size_t)take * 128, g_tcq_dump.fp);
+    g_tcq_dump.written += take;
+    if (g_tcq_dump.written >= g_tcq_dump.max_vecs) {
+        fflush(g_tcq_dump.fp);
+        fprintf(stderr, "TCQ activation-dump: cap reached (%ld vecs) -> closing %s\n", g_tcq_dump.written, "dump file");
+        float * nullp = nullptr;
+        CUDA_CHECK(cudaMemcpyToSymbol(d_tcq_dump_buf, &nullp, sizeof(float*)));
+        if (g_tcq_dump.dev)  { CUDA_CHECK(cudaFree(g_tcq_dump.dev));      g_tcq_dump.dev  = nullptr; g_tcq_dump.dev_vecs  = 0; }
+        if (g_tcq_dump.host) { CUDA_CHECK(cudaFreeHost(g_tcq_dump.host)); g_tcq_dump.host = nullptr; g_tcq_dump.host_vecs = 0; }
+        g_tcq_dump.active = false;
+    }
+}
+
+template<typename idx_t>
+static void set_rows_cuda_turbo3_tcq(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
+    const float * src0_d = (const float *)src0->data;
+    const idx_t * src1_d = (const idx_t *)src1->data;
+    GGML_TENSOR_BINARY_OP_LOCALS
+    GGML_ASSERT(ne00 % QK_TURBO3_TCQ == 0);
+    cudaStream_t stream = ctx.stream();
+    const int64_t ne_total_groups = (ne00 * ne01 * ne02 * ne03) / QK_TURBO3_TCQ;
+    if (ne_total_groups <= 0 || !(ne00 > 0 && ne01 > 0 && ne02 > 0 && ne11 > 0 && ne12 > 0)) return;
+
+    // one-time per-device: optional codebook override (TURBO_TCQ_CB) + adaptive-alpha env
+    static bool tcq3_cb_loaded[GGML_CUDA_MAX_DEVICES] = {};
+    if (!tcq3_cb_loaded[ctx.device]) {
+        tcq3_cb_loaded[ctx.device] = true;
+        const char *cb_path = getenv("TURBO_TCQ_CB");
+        if (cb_path) {
+            float cb[512]; FILE *f = fopen(cb_path, "rb");
+            if (f && fread(cb, sizeof(float), 512, f) == 512) { fclose(f); CUDA_CHECK(cudaMemcpyToSymbol(d_turbo3_tcq_codebook, cb, 512*sizeof(float)));
+                tcq3_set_decode_codebook(cb); // opencoti-hook: TCQ A′ — symmetric (encode+decode) override
+                // opencoti-hook: asymmetric TCQ (#513) — also override the per-TU d_turbo3 copies in the
+                // two asymmetric instance TUs so the runtime codebook reaches asym decode (not just symmetric).
+                tcq3_set_decode_codebook_t3t2(cb); tcq3_set_decode_codebook_t2t3(cb);
+                fprintf(stderr, "TCQ encode+decode: loaded codebook from %s (device %d)\n", cb_path, ctx.device); }
+            else { if (f) fclose(f); fprintf(stderr, "TCQ encode: FAILED to load codebook from %s\n", cb_path); }
+        }
+        load_tcq_norm_alpha(ctx.device);
+    }
+
+    const int64_t s01_f = nb01/sizeof(float), s02_f = nb02/sizeof(float), s03_f = nb03/sizeof(float);
+    const int64_t s10_i = nb10/sizeof(idx_t), s11_i = nb11/sizeof(idx_t), s12_i = nb12/sizeof(idx_t);
+    const int iq_is_k = (strncmp(dst->name, "cache_k_", 8) == 0) ? 1 : 0;
+
+    // one-time per-device: prefer the 8 KB/block shared-memory backtrace (opt-in); else global fallback buffer
+    static int  tcq3_use_shared_bt[GGML_CUDA_MAX_DEVICES] = {};
+    static bool tcq3_bt_checked   [GGML_CUDA_MAX_DEVICES] = {};
+    constexpr int tcq3_bt_shared_bytes = 128 * 64;
+    if (!tcq3_bt_checked[ctx.device]) {
+        tcq3_bt_checked[ctx.device] = true;
+        const char * env = getenv("TURBO_TCQ_SHARED_BT");
+        if (!env || atoi(env) != 0) {
+            int max_shared_optin = 0;
+            CUDA_CHECK(cudaDeviceGetAttribute(&max_shared_optin, cudaDevAttrMaxSharedMemoryPerBlockOptin, ctx.device));
+            if (max_shared_optin >= tcq3_bt_shared_bytes) {
+                if (tcq_fp64_enabled())
+                    CUDA_SET_SHARED_MEMORY_LIMIT((k_set_rows_turbo3_tcq<idx_t, double>), tcq3_bt_shared_bytes);
+                else
+                    CUDA_SET_SHARED_MEMORY_LIMIT(k_set_rows_turbo3_tcq<idx_t>, tcq3_bt_shared_bytes);
+                tcq3_use_shared_bt[ctx.device] = 1;
+            }
+        }
+    }
+    if (!tcq3_use_shared_bt[ctx.device]) ensure_tcq_bt_buf(ctx.device, ne_total_groups * 128 * 64);
+
+    const uint3 ne00_fd = init_fastdiv_values((uint32_t) ne00);
+    const uint3 ne01_fd = init_fastdiv_values((uint32_t) ne01);
+    const uint3 ne02_fd = init_fastdiv_values((uint32_t) ne02);
+    const uint3 ne11_fd = init_fastdiv_values((uint32_t) ne11);
+    const uint3 ne12_fd = init_fastdiv_values((uint32_t) ne12);
+    const int shared_bytes = tcq3_use_shared_bt[ctx.device] ? tcq3_bt_shared_bytes : 0;
+    // opencoti-hook: TCQ activation-dump (F-D, #515) — bind the staging buffer for this launch (no-op unless
+    // TURBO_TCQ_DUMP_ACTS is set). tcq_dump_init() is idempotent; dump_dev is nullptr when inactive/cap-reached,
+    // in which case d_tcq_dump_buf is left at its nullptr default and the kernel's dump branch is a uniform no-op.
+    tcq_dump_init();
+    float * dump_dev = tcq_dump_pre_launch((long) ne_total_groups);
+    if (dump_dev) CUDA_CHECK(cudaMemcpyToSymbol(d_tcq_dump_buf, &dump_dev, sizeof(float*)));
+    if (tcq_fp64_enabled())
+        k_set_rows_turbo3_tcq<idx_t, double><<<(int)ne_total_groups, 512, shared_bytes, stream>>>(
+            src0_d, src1_d, (block_turbo3_tcq *)dst->data,
+            ne_total_groups, tcq_bt_buf[ctx.device], tcq3_use_shared_bt[ctx.device], ne00, ne01, ne02, ne10, ne11, ne12, ne13,
+            s01_f, s02_f, s03_f, s10_i, s11_i, s12_i, iq_is_k, nb1, nb2, nb3,
+            ne00_fd, ne01_fd, ne02_fd, ne11_fd, ne12_fd);
+    else
+        k_set_rows_turbo3_tcq<idx_t><<<(int)ne_total_groups, 512, shared_bytes, stream>>>(
+            src0_d, src1_d, (block_turbo3_tcq *)dst->data,
+            ne_total_groups, tcq_bt_buf[ctx.device], tcq3_use_shared_bt[ctx.device], ne00, ne01, ne02, ne10, ne11, ne12, ne13,
+            s01_f, s02_f, s03_f, s10_i, s11_i, s12_i, iq_is_k, nb1, nb2, nb3,
+            ne00_fd, ne01_fd, ne02_fd, ne11_fd, ne12_fd);
+    if (dump_dev) tcq_dump_post_launch((long) ne_total_groups, stream);
+}
+
+template<typename idx_t>
+static void set_rows_cuda_turbo2_tcq(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
+    const float * src0_d = (const float *)src0->data;
+    const idx_t * src1_d = (const idx_t *)src1->data;
+    GGML_TENSOR_BINARY_OP_LOCALS
+    GGML_ASSERT(ne00 % QK_TURBO2_TCQ == 0);
+    cudaStream_t stream = ctx.stream();
+    const int64_t ne_total_groups = (ne00 * ne01 * ne02 * ne03) / QK_TURBO2_TCQ;
+    if (ne_total_groups <= 0 || !(ne00 > 0 && ne01 > 0 && ne02 > 0 && ne11 > 0 && ne12 > 0)) return;
+
+    static bool tcq2_cb_loaded[GGML_CUDA_MAX_DEVICES] = {};
+    if (!tcq2_cb_loaded[ctx.device]) {
+        tcq2_cb_loaded[ctx.device] = true;
+        const char *cb_path = getenv("TURBO_TCQ_CB2");
+        if (cb_path) {
+            float cb[256]; FILE *f = fopen(cb_path, "rb");
+            if (f && fread(cb, sizeof(float), 256, f) == 256) { fclose(f); CUDA_CHECK(cudaMemcpyToSymbol(d_turbo2_tcq_codebook, cb, 256*sizeof(float)));
+                tcq2_set_decode_codebook(cb); // opencoti-hook: TCQ A′ — symmetric (encode+decode) override
+                // opencoti-hook: asymmetric TCQ (#513) — also override the per-TU d_turbo2 copies in the
+                // two asymmetric instance TUs so the runtime 2-bit codebook reaches asym decode too.
+                tcq2_set_decode_codebook_t3t2(cb); tcq2_set_decode_codebook_t2t3(cb);
+                fprintf(stderr, "TCQ2 encode+decode: loaded 2-bit codebook from %s (device %d)\n", cb_path, ctx.device); }
+            else { if (f) fclose(f); fprintf(stderr, "TCQ2 encode: FAILED to load codebook from %s\n", cb_path); }
+        }
+        load_tcq_norm_alpha(ctx.device);
+    }
+
+    const int64_t s01_f = nb01/sizeof(float), s02_f = nb02/sizeof(float), s03_f = nb03/sizeof(float);
+    const int64_t s10_i = nb10/sizeof(idx_t), s11_i = nb11/sizeof(idx_t), s12_i = nb12/sizeof(idx_t);
+    const int iq_is_k = (strncmp(dst->name, "cache_k_", 8) == 0) ? 1 : 0;
+
+    static int  tcq2_use_shared_bt[GGML_CUDA_MAX_DEVICES] = {};
+    static bool tcq2_bt_checked   [GGML_CUDA_MAX_DEVICES] = {};
+    constexpr int tcq2_bt_shared_bytes = 128 * 64;
+    if (!tcq2_bt_checked[ctx.device]) {
+        tcq2_bt_checked[ctx.device] = true;
+        const char * env = getenv("TURBO_TCQ_SHARED_BT");
+        if (!env || atoi(env) != 0) {
+            int max_shared_optin = 0;
+            CUDA_CHECK(cudaDeviceGetAttribute(&max_shared_optin, cudaDevAttrMaxSharedMemoryPerBlockOptin, ctx.device));
+            if (max_shared_optin >= tcq2_bt_shared_bytes) {
+                if (tcq_fp64_enabled())
+                    CUDA_SET_SHARED_MEMORY_LIMIT((k_set_rows_turbo2_tcq<idx_t, double>), tcq2_bt_shared_bytes);
+                else
+                    CUDA_SET_SHARED_MEMORY_LIMIT(k_set_rows_turbo2_tcq<idx_t>, tcq2_bt_shared_bytes);
+                tcq2_use_shared_bt[ctx.device] = 1;
+            }
+        }
+    }
+    if (!tcq2_use_shared_bt[ctx.device]) ensure_tcq_bt_buf(ctx.device, ne_total_groups * 128 * 64);
+
+    const uint3 ne00_fd = init_fastdiv_values((uint32_t) ne00);
+    const uint3 ne01_fd = init_fastdiv_values((uint32_t) ne01);
+    const uint3 ne02_fd = init_fastdiv_values((uint32_t) ne02);
+    const uint3 ne11_fd = init_fastdiv_values((uint32_t) ne11);
+    const uint3 ne12_fd = init_fastdiv_values((uint32_t) ne12);
+    const int shared_bytes = tcq2_use_shared_bt[ctx.device] ? tcq2_bt_shared_bytes : 0;
+    if (tcq_fp64_enabled())
+        k_set_rows_turbo2_tcq<idx_t, double><<<(int)ne_total_groups, 256, shared_bytes, stream>>>(
+            src0_d, src1_d, (block_turbo2_tcq *)dst->data,
+            ne_total_groups, tcq_bt_buf[ctx.device], tcq2_use_shared_bt[ctx.device], ne00, ne01, ne02, ne10, ne11, ne12, ne13,
+            s01_f, s02_f, s03_f, s10_i, s11_i, s12_i, iq_is_k, nb1, nb2, nb3,
+            ne00_fd, ne01_fd, ne02_fd, ne11_fd, ne12_fd);
+    else
+        k_set_rows_turbo2_tcq<idx_t><<<(int)ne_total_groups, 256, shared_bytes, stream>>>(
+            src0_d, src1_d, (block_turbo2_tcq *)dst->data,
+            ne_total_groups, tcq_bt_buf[ctx.device], tcq2_use_shared_bt[ctx.device], ne00, ne01, ne02, ne10, ne11, ne12, ne13,
+            s01_f, s02_f, s03_f, s10_i, s11_i, s12_i, iq_is_k, nb1, nb2, nb3,
+            ne00_fd, ne01_fd, ne02_fd, ne11_fd, ne12_fd);
+}
+
 template<typename idx_t>
 static void set_rows_cuda_turbo3(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
     const float * src0_d = (const float *)src0->data;
@@ -1094,6 +1363,18 @@ static void set_rows_cuda(ggml_backend_cuda_context & ctx, const ggml_tensor * s
             nb1, nb2, nb3,
             stream
         );
+    // opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port) — KV-write via the generic quantized set_rows
+    // template (mirrors q5_0); quantize_f32_q6_0_block is in cpy-utils.cuh. See docs/features/poly_kv.md.
+    } else if (dst->type == GGML_TYPE_Q6_0) {
+        set_rows_cuda_quant<idx_t, block_q6_0, QK6_0, quantize_f32_q6_0_block>(
+            src0_d, src1_d, (block_q6_0*)dst->data,
+            ne00, ne01, ne02, ne03,
+            ne10, ne11, ne12, ne13,
+            nb01, nb02, nb03,
+            nb10, nb11, nb12,
+            nb1, nb2, nb3,
+            stream
+        );
     } else if (dst->type == GGML_TYPE_Q8_0) {
         set_rows_cuda_quant<idx_t, block_q8_0, QK8_0, quantize_f32_q8_0_block>(
             src0_d, src1_d, (block_q8_0*)dst->data,
@@ -1126,6 +1407,12 @@ static void set_rows_cuda(ggml_backend_cuda_context & ctx, const ggml_tensor * s
     } else if (dst->type == GGML_TYPE_TURBO8_0) {
         // opencoti-hook: turboquant tier family — see docs/features/poly_kv.md (M6-S2)
         set_rows_cuda_turbo8<idx_t>(ctx, src0, src1, dst);
+    } else if (dst->type == GGML_TYPE_TURBO3_TCQ) {
+        // opencoti-hook: TCQ (F5 M6 C5) — Viterbi 3-bit trellis encode (turbo-tcq-cuda.cuh)
+        set_rows_cuda_turbo3_tcq<idx_t>(ctx, src0, src1, dst);
+    } else if (dst->type == GGML_TYPE_TURBO2_TCQ) {
+        // opencoti-hook: TCQ (F5 M6 C5) — Viterbi 2-bit trellis encode (turbo-tcq-cuda.cuh)
+        set_rows_cuda_turbo2_tcq<idx_t>(ctx, src0, src1, dst);
     } else {
         GGML_ABORT("unsupported type %s", ggml_type_name(dst->type));
     }
diff --git a/llama.cpp/ggml/src/ggml-cuda/tcq-decode-cb.cuh b/llama.cpp/ggml/src/ggml-cuda/tcq-decode-cb.cuh
new file mode 100644
index 000000000..5819745b9
--- /dev/null
+++ b/llama.cpp/ggml/src/ggml-cuda/tcq-decode-cb.cuh
@@ -0,0 +1,24 @@
+#pragma once
+// opencoti-hook: TCQ A′ symmetric codebook override — see docs/evaluations/context.md §5a-TCQ.
+//
+// The decode-side TCQ codebooks (d_turbo{3,2}_tcq_codebook_fattn) live as `static __constant__` in
+// fattn-tcq.cuh and are replicated PER-TU (the CUDA DSO compiles each .cu separately, no -rdc). The
+// kernel that READS each copy is the FA-VEC decode instantiated in
+// template-instances/fattn-vec-instance-turbo{3,2}_tcq-*.cu, so the only copy that matters is the one
+// in THAT translation unit. These setters are therefore DEFINED in those instance TUs (where the live
+// copy + cudaMemcpyToSymbol can reach it) and DECLARED here for set-rows.cu's TURBO_TCQ_CB / CB2 loader
+// to call. Without this, TURBO_TCQ_CB overrides only the ENCODE symbol and FA decode keeps the baked
+// codebook → mismatch corruption (footgun A′). With it, a runtime codebook override is symmetric.
+void tcq3_set_decode_codebook(const float * cb);   // 512 floats (turbo3_tcq)
+void tcq2_set_decode_codebook(const float * cb);   // 256 floats (turbo2_tcq)
+
+// opencoti-hook: asymmetric TCQ (#513). The asymmetric instance TUs (turbo3_tcq/turbo2_tcq + reverse)
+// each own their OWN per-TU copies of BOTH d_turbo3 + d_turbo2 codebooks (the asym decode reads both: K
+// from one tier, V from the other), so each TU needs its own setter pair — uniquely named to avoid
+// duplicate-symbol link errors with the symmetric setters above. set-rows.cu calls all of these in the
+// TURBO_TCQ_CB / CB2 apply so a runtime override reaches the asymmetric decode TUs too (else they keep
+// the baked codebook while the symmetric path uses the override → inconsistent KLD).
+void tcq3_set_decode_codebook_t3t2(const float * cb);  // turbo3_tcq-K / turbo2_tcq-V TU: its d_turbo3 (512)
+void tcq2_set_decode_codebook_t3t2(const float * cb);  // turbo3_tcq-K / turbo2_tcq-V TU: its d_turbo2 (256)
+void tcq3_set_decode_codebook_t2t3(const float * cb);  // turbo2_tcq-K / turbo3_tcq-V TU: its d_turbo3 (512)
+void tcq2_set_decode_codebook_t2t3(const float * cb);  // turbo2_tcq-K / turbo3_tcq-V TU: its d_turbo2 (256)
diff --git a/llama.cpp/ggml/src/ggml-cuda/template-instances/fattn-vec-instance-f16-q6_0.cu b/llama.cpp/ggml/src/ggml-cuda/template-instances/fattn-vec-instance-f16-q6_0.cu
new file mode 100644
index 000000000..1e81a509a
--- /dev/null
+++ b/llama.cpp/ggml/src/ggml-cuda/template-instances/fattn-vec-instance-f16-q6_0.cu
@@ -0,0 +1,7 @@
+// opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port). Mirrors fattn-vec-instance-f16-q5_0.cu.
+
+#include "../fattn-vec.cuh"
+
+DECL_FATTN_VEC_CASE( 64, GGML_TYPE_F16, GGML_TYPE_Q6_0);
+DECL_FATTN_VEC_CASE(128, GGML_TYPE_F16, GGML_TYPE_Q6_0);
+DECL_FATTN_VEC_CASE(256, GGML_TYPE_F16, GGML_TYPE_Q6_0);
diff --git a/llama.cpp/ggml/src/ggml-cuda/template-instances/fattn-vec-instance-q6_0-f16.cu b/llama.cpp/ggml/src/ggml-cuda/template-instances/fattn-vec-instance-q6_0-f16.cu
new file mode 100644
index 000000000..0f2c5b89c
--- /dev/null
+++ b/llama.cpp/ggml/src/ggml-cuda/template-instances/fattn-vec-instance-q6_0-f16.cu
@@ -0,0 +1,7 @@
+// opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port). Mirrors fattn-vec-instance-q5_0-f16.cu.
+
+#include "../fattn-vec.cuh"
+
+DECL_FATTN_VEC_CASE( 64, GGML_TYPE_Q6_0, GGML_TYPE_F16);
+DECL_FATTN_VEC_CASE(128, GGML_TYPE_Q6_0, GGML_TYPE_F16);
+DECL_FATTN_VEC_CASE(256, GGML_TYPE_Q6_0, GGML_TYPE_F16);
diff --git a/llama.cpp/ggml/src/ggml-cuda/template-instances/fattn-vec-instance-q6_0-q5_0.cu b/llama.cpp/ggml/src/ggml-cuda/template-instances/fattn-vec-instance-q6_0-q5_0.cu
new file mode 100644
index 000000000..93a414520
--- /dev/null
+++ b/llama.cpp/ggml/src/ggml-cuda/template-instances/fattn-vec-instance-q6_0-q5_0.cu
@@ -0,0 +1,7 @@
+// opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port). Mirrors fattn-vec-instance-q5_0-q4_0.cu form.
+
+#include "../fattn-vec.cuh"
+
+DECL_FATTN_VEC_CASE( 64, GGML_TYPE_Q6_0, GGML_TYPE_Q5_0);
+DECL_FATTN_VEC_CASE(128, GGML_TYPE_Q6_0, GGML_TYPE_Q5_0);
+DECL_FATTN_VEC_CASE(256, GGML_TYPE_Q6_0, GGML_TYPE_Q5_0);
diff --git a/llama.cpp/ggml/src/ggml-cuda/template-instances/fattn-vec-instance-q6_0-q6_0.cu b/llama.cpp/ggml/src/ggml-cuda/template-instances/fattn-vec-instance-q6_0-q6_0.cu
new file mode 100644
index 000000000..04ac8a57d
--- /dev/null
+++ b/llama.cpp/ggml/src/ggml-cuda/template-instances/fattn-vec-instance-q6_0-q6_0.cu
@@ -0,0 +1,7 @@
+// opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port). Mirrors fattn-vec-instance-q5_0-q5_0.cu.
+
+#include "../fattn-vec.cuh"
+
+DECL_FATTN_VEC_CASE( 64, GGML_TYPE_Q6_0, GGML_TYPE_Q6_0);
+DECL_FATTN_VEC_CASE(128, GGML_TYPE_Q6_0, GGML_TYPE_Q6_0);
+DECL_FATTN_VEC_CASE(256, GGML_TYPE_Q6_0, GGML_TYPE_Q6_0);
diff --git a/llama.cpp/ggml/src/ggml-cuda/template-instances/fattn-vec-instance-q8_0-q6_0.cu b/llama.cpp/ggml/src/ggml-cuda/template-instances/fattn-vec-instance-q8_0-q6_0.cu
new file mode 100644
index 000000000..563f54cc4
--- /dev/null
+++ b/llama.cpp/ggml/src/ggml-cuda/template-instances/fattn-vec-instance-q8_0-q6_0.cu
@@ -0,0 +1,7 @@
+// opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port). Mirrors fattn-vec-instance-q8_0-q5_0.cu.
+
+#include "../fattn-vec.cuh"
+
+DECL_FATTN_VEC_CASE( 64, GGML_TYPE_Q8_0, GGML_TYPE_Q6_0);
+DECL_FATTN_VEC_CASE(128, GGML_TYPE_Q8_0, GGML_TYPE_Q6_0);
+DECL_FATTN_VEC_CASE(256, GGML_TYPE_Q8_0, GGML_TYPE_Q6_0);
diff --git a/llama.cpp/ggml/src/ggml-cuda/template-instances/fattn-vec-instance-turbo2_tcq-turbo2_tcq.cu b/llama.cpp/ggml/src/ggml-cuda/template-instances/fattn-vec-instance-turbo2_tcq-turbo2_tcq.cu
new file mode 100644
index 000000000..9f4e7fbc1
--- /dev/null
+++ b/llama.cpp/ggml/src/ggml-cuda/template-instances/fattn-vec-instance-turbo2_tcq-turbo2_tcq.cu
@@ -0,0 +1,17 @@
+// This file mirrors the autogenerated turbo instance files; do not edit manually.
+// opencoti-hook: TCQ (trellis-coded quant) FA-VEC decode — F5 M6 C5 (#447/#500). See UPSTREAM_SYNC.md.
+
+#include "../fattn-vec.cuh"
+
+DECL_FATTN_VEC_CASE( 64, GGML_TYPE_TURBO2_TCQ, GGML_TYPE_TURBO2_TCQ);
+DECL_FATTN_VEC_CASE(128, GGML_TYPE_TURBO2_TCQ, GGML_TYPE_TURBO2_TCQ);
+DECL_FATTN_VEC_CASE(256, GGML_TYPE_TURBO2_TCQ, GGML_TYPE_TURBO2_TCQ);
+DECL_FATTN_VEC_CASE(512, GGML_TYPE_TURBO2_TCQ, GGML_TYPE_TURBO2_TCQ); // opencoti TCQ: Gemma-4 GLOBAL head_dim 512
+
+// opencoti-hook: TCQ A′ symmetric codebook override (decode side) — see tcq-decode-cb.cuh.
+// Defined HERE because d_turbo2_tcq_codebook_fattn is static __constant__ per-TU (this TU owns the
+// live copy the FA-VEC decode above reads).
+#include "../tcq-decode-cb.cuh"
+void tcq2_set_decode_codebook(const float * cb) {
+    CUDA_CHECK(cudaMemcpyToSymbol(d_turbo2_tcq_codebook_fattn, cb, 256 * sizeof(float)));
+}
diff --git a/llama.cpp/ggml/src/ggml-cuda/template-instances/fattn-vec-instance-turbo2_tcq-turbo3_tcq.cu b/llama.cpp/ggml/src/ggml-cuda/template-instances/fattn-vec-instance-turbo2_tcq-turbo3_tcq.cu
new file mode 100644
index 000000000..b51a16283
--- /dev/null
+++ b/llama.cpp/ggml/src/ggml-cuda/template-instances/fattn-vec-instance-turbo2_tcq-turbo3_tcq.cu
@@ -0,0 +1,22 @@
+// This file mirrors the autogenerated turbo instance files; do not edit manually.
+// opencoti-hook: asymmetric TCQ (#513) FA-VEC decode — turbo2_tcq K / turbo3_tcq V (strong-V).
+// See UPSTREAM_SYNC.md and docs/evaluations/context.md §5a-TCQ.
+
+#include "../fattn-vec.cuh"
+
+DECL_FATTN_VEC_CASE( 64, GGML_TYPE_TURBO2_TCQ, GGML_TYPE_TURBO3_TCQ);
+DECL_FATTN_VEC_CASE(128, GGML_TYPE_TURBO2_TCQ, GGML_TYPE_TURBO3_TCQ);
+DECL_FATTN_VEC_CASE(256, GGML_TYPE_TURBO2_TCQ, GGML_TYPE_TURBO3_TCQ);
+DECL_FATTN_VEC_CASE(512, GGML_TYPE_TURBO2_TCQ, GGML_TYPE_TURBO3_TCQ); // opencoti TCQ: Gemma-4 GLOBAL head_dim 512
+
+// opencoti-hook: asymmetric TCQ A′ codebook override (decode side) — this TU owns its OWN per-TU static
+// __constant__ d_turbo3_tcq_codebook_fattn (V decode) AND d_turbo2_tcq_codebook_fattn (K decode), so BOTH
+// need their own setter here (uniquely named vs the symmetric TUs to avoid duplicate symbols). set-rows.cu
+// calls these from the TURBO_TCQ_CB / CB2 loader so a runtime override reaches this TU's decode too.
+#include "../tcq-decode-cb.cuh"
+void tcq3_set_decode_codebook_t2t3(const float * cb) {
+    CUDA_CHECK(cudaMemcpyToSymbol(d_turbo3_tcq_codebook_fattn, cb, 512 * sizeof(float)));
+}
+void tcq2_set_decode_codebook_t2t3(const float * cb) {
+    CUDA_CHECK(cudaMemcpyToSymbol(d_turbo2_tcq_codebook_fattn, cb, 256 * sizeof(float)));
+}
diff --git a/llama.cpp/ggml/src/ggml-cuda/template-instances/fattn-vec-instance-turbo3_tcq-turbo2_tcq.cu b/llama.cpp/ggml/src/ggml-cuda/template-instances/fattn-vec-instance-turbo3_tcq-turbo2_tcq.cu
new file mode 100644
index 000000000..da175f9c1
--- /dev/null
+++ b/llama.cpp/ggml/src/ggml-cuda/template-instances/fattn-vec-instance-turbo3_tcq-turbo2_tcq.cu
@@ -0,0 +1,22 @@
+// This file mirrors the autogenerated turbo instance files; do not edit manually.
+// opencoti-hook: asymmetric TCQ (#513) FA-VEC decode — turbo3_tcq K / turbo2_tcq V (strong-K).
+// See UPSTREAM_SYNC.md and docs/evaluations/context.md §5a-TCQ.
+
+#include "../fattn-vec.cuh"
+
+DECL_FATTN_VEC_CASE( 64, GGML_TYPE_TURBO3_TCQ, GGML_TYPE_TURBO2_TCQ);
+DECL_FATTN_VEC_CASE(128, GGML_TYPE_TURBO3_TCQ, GGML_TYPE_TURBO2_TCQ);
+DECL_FATTN_VEC_CASE(256, GGML_TYPE_TURBO3_TCQ, GGML_TYPE_TURBO2_TCQ);
+DECL_FATTN_VEC_CASE(512, GGML_TYPE_TURBO3_TCQ, GGML_TYPE_TURBO2_TCQ); // opencoti TCQ: Gemma-4 GLOBAL head_dim 512
+
+// opencoti-hook: asymmetric TCQ A′ codebook override (decode side) — this TU owns its OWN per-TU static
+// __constant__ d_turbo3_tcq_codebook_fattn (K decode) AND d_turbo2_tcq_codebook_fattn (V decode), so BOTH
+// need their own setter here (uniquely named vs the symmetric TUs to avoid duplicate symbols). set-rows.cu
+// calls these from the TURBO_TCQ_CB / CB2 loader so a runtime override reaches this TU's decode too.
+#include "../tcq-decode-cb.cuh"
+void tcq3_set_decode_codebook_t3t2(const float * cb) {
+    CUDA_CHECK(cudaMemcpyToSymbol(d_turbo3_tcq_codebook_fattn, cb, 512 * sizeof(float)));
+}
+void tcq2_set_decode_codebook_t3t2(const float * cb) {
+    CUDA_CHECK(cudaMemcpyToSymbol(d_turbo2_tcq_codebook_fattn, cb, 256 * sizeof(float)));
+}
diff --git a/llama.cpp/ggml/src/ggml-cuda/template-instances/fattn-vec-instance-turbo3_tcq-turbo3_tcq.cu b/llama.cpp/ggml/src/ggml-cuda/template-instances/fattn-vec-instance-turbo3_tcq-turbo3_tcq.cu
new file mode 100644
index 000000000..906910dde
--- /dev/null
+++ b/llama.cpp/ggml/src/ggml-cuda/template-instances/fattn-vec-instance-turbo3_tcq-turbo3_tcq.cu
@@ -0,0 +1,17 @@
+// This file mirrors the autogenerated turbo instance files; do not edit manually.
+// opencoti-hook: TCQ (trellis-coded quant) FA-VEC decode — F5 M6 C5 (#447/#500). See UPSTREAM_SYNC.md.
+
+#include "../fattn-vec.cuh"
+
+DECL_FATTN_VEC_CASE( 64, GGML_TYPE_TURBO3_TCQ, GGML_TYPE_TURBO3_TCQ);
+DECL_FATTN_VEC_CASE(128, GGML_TYPE_TURBO3_TCQ, GGML_TYPE_TURBO3_TCQ);
+DECL_FATTN_VEC_CASE(256, GGML_TYPE_TURBO3_TCQ, GGML_TYPE_TURBO3_TCQ);
+DECL_FATTN_VEC_CASE(512, GGML_TYPE_TURBO3_TCQ, GGML_TYPE_TURBO3_TCQ); // opencoti TCQ: Gemma-4 GLOBAL head_dim 512
+
+// opencoti-hook: TCQ A′ symmetric codebook override (decode side) — see tcq-decode-cb.cuh.
+// Defined HERE because d_turbo3_tcq_codebook_fattn is static __constant__ per-TU and the live copy the
+// FA-VEC decode above reads is this TU's. cudaMemcpyToSymbol on a static symbol only works in-TU.
+#include "../tcq-decode-cb.cuh"
+void tcq3_set_decode_codebook(const float * cb) {
+    CUDA_CHECK(cudaMemcpyToSymbol(d_turbo3_tcq_codebook_fattn, cb, 512 * sizeof(float)));
+}
diff --git a/llama.cpp/ggml/src/ggml-cuda/turbo-tcq-cuda.cuh b/llama.cpp/ggml/src/ggml-cuda/turbo-tcq-cuda.cuh
new file mode 100644
index 000000000..f5b3ebba7
--- /dev/null
+++ b/llama.cpp/ggml/src/ggml-cuda/turbo-tcq-cuda.cuh
@@ -0,0 +1,747 @@
+#pragma once
+// opencoti-hook: TCQ (F5 M6 C5) — Trellis-Coded Quantization KV ENCODE (CUDA).
+// Ported from buun (github.com/spiritbuun/buun-llama-cpp, ggml-cuda/turbo-quant-cuda.cuh):
+// the 512/256-state right-shift Viterbi encode kernels, the GET_ROWS dequant helpers, and the
+// trained compiled-in codebooks (credit spiritbuun). Adaptations vs buun (all additive/syncable):
+//   * reuses our TURBO3_WHT_SIGNS1/2 (turbo-dequant.cuh) — verified byte-identical to buun's
+//     d_turbo_wht_signs1/2, so the trained codebook stays valid;
+//   * InnerQ calibrate + per-channel scale blocks STRIPPED (identity) — matches our turbo3 encode
+//     ("InnerQ stub-disabled"); avoids colliding with our head_dim-aware d_innerq_* in set-rows.cu.
+//     TCQ x InnerQ composition is a deliberate later A/B (our InnerQ is head_dim-aware, buun's 128-only);
+//   * diagnostic dump blocks (d_tcq_dump_*) and turbo_extract_append removed;
+//   * fp32 baseline (byte-faithful to buun). fp64-selectable encode (env TURBO_TCQ_FP64) lands as a
+//     follow-on once the fp32 trellis is correctness-gated (logit-equiv + KLD vs buun).
+// Provenance + design: docs/evaluations/context.md §5.
+#include "common.cuh"
+#include "turbo-dequant.cuh"   // TURBO3_WHT_SIGNS1/2 (reused; byte-identical to buun)
+
+// NOTE: the context-adaptive norm-scaling symbols d_tcq_norm_alpha (K) / d_tcq_norm_alpha_v (V)
+// are declared in the extracted buun section below (defaults 1.0 / 1.04). Stage-4 tunes them via
+// cudaMemcpyToSymbol from set-rows.cu (TURBO_TCQ_ALPHA / _ALPHA_V env).
+
+static __constant__ float d_turbo3_tcq_codebook[512] = {
+    -0.14559399f, -0.09062801f, -0.054925077f, -0.03699251f, -0.006363985f, +0.026264573f, +0.067378916f, +0.121981815f,
+    -0.18648055f, -0.106522456f, -0.052047577f, -0.011695214f, +0.021953275f, +0.059698727f, +0.09831437f, +0.16083933f,
+    -0.16390342f, -0.12639847f, -0.09513180f, -0.05938352f, -0.028396897f, +0.005973862f, +0.049104784f, +0.11334257f,
+    -0.25952467f, -0.079778515f, -0.036024813f, +0.0003641268f, +0.031858794f, +0.073280424f, +0.11835553f, +0.19738495f,
+    -0.14218009f, -0.10224814f, -0.062498566f, -0.027066832f, +0.00393002f, +0.04069300f, +0.08257346f, +0.14548601f,
+    -0.18673635f, -0.13438253f, -0.088401966f, -0.05205436f, -0.02032501f, +0.012399545f, +0.05127183f, +0.10316186f,
+    -0.10807011f, -0.065903045f, -0.032206114f, -0.0062006037f, +0.020679146f, +0.04422085f, +0.08313074f, +0.16821936f,
+    -0.22979105f, -0.14431947f, -0.07689272f, -0.02755307f, +0.009225173f, +0.046684854f, +0.08834142f, +0.13766693f,
+    -0.22114082f, -0.12612148f, -0.06890522f, -0.016128855f, +0.03691900f, +0.08474852f, +0.14940020f, +0.23229980f,
+    -0.14933491f, -0.099693604f, -0.06738499f, -0.037100967f, -0.009332986f, +0.023535024f, +0.060272533f, +0.109464675f,
+    -0.20200425f, -0.07398328f, -0.038700905f, -0.01714807f, +0.011161969f, +0.04528101f, +0.08902637f, +0.19573534f,
+    -0.16645233f, -0.124482535f, -0.089342155f, -0.04427387f, -0.007353691f, +0.028033108f, +0.066108435f, +0.15552913f,
+    -0.22295763f, -0.059887577f, -0.018804537f, +0.020141022f, +0.059682943f, +0.097920544f, +0.14080113f, +0.25698325f,
+    -0.14248224f, -0.089685425f, -0.050101686f, -0.017257255f, +0.011412255f, +0.040830314f, +0.07400172f, +0.11997315f,
+    -0.18649384f, -0.113997504f, -0.067775466f, -0.033394672f, +0.006586988f, +0.05312057f, +0.10433043f, +0.22344802f,
+    -0.16138338f, -0.108194515f, -0.07600300f, -0.05135381f, -0.023365447f, +0.0087320795f, +0.045431953f, +0.09113002f,
+    -0.12630440f, -0.07225349f, -0.032280035f, +0.0029231994f, +0.019239848f, +0.05081419f, +0.077840395f, +0.121695265f,
+    -0.08928155f, -0.044983763f, -0.009889568f, +0.020831043f, +0.05684458f, +0.09409702f, +0.13867535f, +0.19084482f,
+    -0.14182915f, -0.11380146f, -0.06904074f, -0.002002765f, +0.034864165f, +0.070399575f, +0.11403063f, +0.15394832f,
+    -0.10876417f, -0.056122433f, -0.02267638f, +0.011113975f, +0.039639056f, +0.074084364f, +0.10155376f, +0.12540291f,
+    -0.17693359f, -0.13940524f, -0.10049578f, -0.06796275f, -0.036915872f, +0.00062823476f, +0.042142134f, +0.17906062f,
+    -0.09253492f, -0.04290128f, -0.006311852f, +0.023908244f, +0.049849935f, +0.078770354f, +0.10818172f, +0.15166481f,
+    -0.12429565f, -0.07392063f, -0.029114135f, +0.0059440783f, +0.042675965f, +0.08425635f, +0.13836108f, +0.18634140f,
+    -0.11795639f, -0.07033707f, -0.034163877f, -0.0008773357f, +0.03334606f, +0.07188203f, +0.12216825f, +0.17097956f,
+    -0.18718453f, -0.14090346f, -0.097799584f, -0.059522875f, -0.019208657f, +0.03079176f, +0.09334672f, +0.15811224f,
+    -0.27198875f, -0.16546582f, -0.11433405f, -0.06933013f, -0.04026183f, -0.0061146915f, +0.029263576f, +0.07322499f,
+    -0.18471734f, -0.102074504f, -0.06492570f, -0.034418534f, -0.009636157f, +0.023043344f, +0.05751496f, +0.09905984f,
+    -0.22826399f, -0.15946552f, -0.09913176f, -0.06585259f, -0.03252090f, +0.001313243f, +0.03556729f, +0.21612854f,
+    -0.13243781f, -0.087299444f, -0.049820945f, -0.016216082f, +0.01799807f, +0.057916876f, +0.09001349f, +0.13221787f,
+    -0.19516511f, -0.120894566f, -0.076130204f, -0.051442243f, -0.029535033f, -0.0020043184f, +0.029452588f, +0.075566076f,
+    -0.27272871f, -0.15841717f, -0.105432935f, -0.06792948f, -0.024532158f, +0.014960791f, +0.054415092f, +0.101517834f,
+    -0.21153601f, -0.15015371f, -0.08676790f, -0.04414934f, -0.0042129597f, +0.033762872f, +0.07589151f, +0.12768789f,
+    -0.090428725f, -0.037582967f, +0.0013173596f, +0.03900247f, +0.06840049f, +0.116906695f, +0.16584939f, +0.25382105f,
+    -0.13446195f, -0.07865091f, -0.039625354f, -0.0028398742f, +0.03019514f, +0.06799379f, +0.11850997f, +0.17521496f,
+    -0.11350345f, -0.058599845f, -0.017512511f, +0.019431496f, +0.055897832f, +0.093173414f, +0.14820710f, +0.22092152f,
+    -0.15165758f, -0.08869354f, -0.04974287f, -0.01705474f, +0.013134752f, +0.04367713f, +0.07733791f, +0.12430801f,
+    -0.09329869f, -0.04673005f, -0.00045857552f, +0.042781368f, +0.07802363f, +0.11887439f, +0.16250038f, +0.28612965f,
+    -0.12571070f, -0.07786012f, -0.03843933f, -0.0075433915f, +0.025822964f, +0.066053316f, +0.12021536f, +0.18341768f,
+    -0.16079275f, -0.04921760f, -0.006114644f, +0.026215268f, +0.05699377f, +0.09813471f, +0.16080129f, +0.23786584f,
+    -0.09980837f, -0.048535258f, -0.0096120685f, +0.025387142f, +0.05979822f, +0.09875251f, +0.14474337f, +0.20324114f,
+    -0.15846540f, -0.09938028f, -0.061492465f, -0.03523542f, -0.0061364113f, +0.024916094f, +0.06037314f, +0.106796466f,
+    -0.20557843f, -0.123237535f, -0.07734871f, -0.044549115f, -0.017114898f, +0.01616654f, +0.049574375f, +0.092319444f,
+    -0.19221115f, -0.14642999f, -0.091701314f, -0.055265956f, -0.021026207f, +0.017720066f, +0.05786183f, +0.110154524f,
+    -0.09956386f, -0.03870283f, +0.003052007f, +0.034851722f, +0.06256365f, +0.09628840f, +0.13979156f, +0.16582295f,
+    -0.18026546f, -0.12448310f, -0.07424377f, -0.03954519f, -0.01221123f, +0.028641058f, +0.100819774f, +0.18240699f,
+    -0.21520759f, -0.15573645f, -0.09820838f, -0.051450998f, -0.012993679f, +0.021135861f, +0.058727216f, +0.105848536f,
+    -0.11207385f, -0.08335689f, -0.048542723f, -0.023198519f, +0.0039304253f, +0.037778318f, +0.07813917f, +0.13106476f,
+    -0.17849164f, -0.120988995f, -0.078016765f, -0.043093704f, -0.016565649f, +0.015182641f, +0.050754096f, +0.09595712f,
+    -0.22132620f, -0.13407415f, -0.065785654f, -0.013291034f, +0.032098345f, +0.07478225f, +0.12431934f, +0.19174045f,
+    -0.095454164f, -0.051898945f, -0.015116375f, -0.012596778f, +0.018636847f, +0.05006925f, +0.087654814f, +0.13754296f,
+    -0.15254061f, -0.09576059f, -0.052086458f, -0.01596074f, +0.017607626f, +0.04778498f, +0.08950204f, +0.14901252f,
+    -0.26057002f, -0.12472382f, -0.074396215f, -0.03764066f, +0.0011168446f, +0.061569117f, +0.10793752f, +0.19771695f,
+    -0.08661132f, -0.045195263f, -0.016098704f, +0.012780116f, +0.040476497f, +0.074102715f, +0.074102715f, +0.12635531f,
+    -0.14047913f, -0.059587404f, -0.016261123f, +0.019801628f, +0.053541403f, +0.096650146f, +0.15005490f, +0.21051759f,
+    -0.22986396f, -0.11964334f, -0.07266585f, -0.026522418f, +0.018169926f, +0.058630653f, +0.100647695f, +0.15919648f,
+    -0.13251697f, -0.077567816f, -0.042766172f, -0.011389967f, +0.01831755f, +0.05304656f, +0.09620367f, +0.15567583f,
+    -0.119819686f, -0.06772876f, -0.028123451f, +0.00876240f, +0.014405836f, +0.048829112f, +0.08422175f, +0.13823749f,
+    -0.16379014f, -0.08956941f, -0.041652776f, +0.008921398f, +0.05473602f, +0.10037984f, +0.16022855f, +0.23457925f,
+    -0.115844205f, -0.05939626f, -0.020390417f, +0.01374377f, +0.044976473f, +0.07873563f, +0.12207942f, +0.18412720f,
+    -0.19048831f, -0.07587487f, -0.03220580f, -0.00011795067f, +0.02721784f, +0.04380719f, +0.07886723f, +0.13193911f,
+    -0.13935551f, -0.092902906f, -0.052706074f, -0.017797327f, +0.015312965f, +0.056098964f, +0.11203423f, +0.24448302f,
+    -0.17986591f, -0.10738580f, -0.06376371f, -0.026595421f, +0.00842492f, +0.04272362f, +0.08608052f, +0.15240218f,
+    -0.10953678f, -0.057022586f, -0.012483291f, +0.024463262f, +0.06076792f, +0.09776234f, +0.12983681f, +0.18648379f,
+    -0.16471463f, -0.089491285f, -0.037574016f, +0.004444791f, +0.039293647f, +0.07845859f, +0.12893885f, +0.23508036f
+};
+
+// TCQ norm alpha: K alpha = 1.0 (no scaling), V alpha = 1.04 (optimal at 2K context).
+// Override via TURBO_TCQ_ALPHA (K) and TURBO_TCQ_ALPHA_V (V) env vars.
+// Alpha is applied at encode time (baked into fp16 norm) — this outperforms decode-time application.
+static __constant__ float d_tcq_norm_alpha = 1.0f;
+static __constant__ float d_tcq_norm_alpha_v = 1.04f;
+
+// opencoti-hook: TCQ activation-dump (F-D, plan lexical-gliding-flurry / #515) — see docs/evaluations/context.md §5a-TCQ.
+// When non-null, k_set_rows_turbo3_tcq stages each group's post-FWHT unit-norm 128-vec (the exact codebook-domain
+// signal the trellis quantizes against d_turbo3_tcq_codebook) into this device buffer at [group*128 + i]. Set
+// per-launch from set-rows.cu ONLY when TURBO_TCQ_DUMP_ACTS is set; stays nullptr otherwise (uniform branch, zero
+// overhead). Feeds F-E's real-27B codebook retrain (the 512-entry turbo3_tcq codebook); the whitened unit-norm
+// vec is codebook-resolution-independent (FWHT happens before the trellis).
+static __constant__ float * d_tcq_dump_buf = nullptr;
+
+// fp64-selectable encode (env TURBO_TCQ_FP64): the Viterbi forward pass and the
+// norm/whiten reductions run in acc_t. acc_t defaults to float — every existing
+// <idx_t> launch resolves to the float instantiation, byte-identical to buun.
+// The host dispatches the <idx_t,double> instantiation when TURBO_TCQ_FP64 is set.
+// Overloaded (not specialized) so the float overload lowers to the exact same
+// sqrtf SASS as the original code → fp32 byte-identity is guaranteed by construction.
+static __device__ __forceinline__ float  tcq_acc_sqrt(float  v) { return sqrtf(v); }
+static __device__ __forceinline__ double tcq_acc_sqrt(double v) { return sqrt(v);  }
+
+// TCQ SET_ROWS encode: Viterbi optimal path with right-shift trellis
+// 512 threads per block (one per trellis state), one block per 128-element group
+// Double-buffered cost arrays + global memory backtrace (128 syncs/group, was 384)
+template<typename idx_t, typename acc_t = float>
+static __global__ void __launch_bounds__(512, 1) k_set_rows_turbo3_tcq(
+        const float * __restrict__ src0, const idx_t * __restrict__ src1,
+        block_turbo3_tcq * __restrict__ dst, const int64_t ne_total_groups,
+        uint8_t * __restrict__ bt_buf,
+        const int use_shared_bt,
+        const int64_t ne00, const int64_t ne01, const int64_t ne02,
+        const int64_t ne10, const int64_t ne11, const int64_t ne12, const int64_t ne13,
+        const int64_t s01, const int64_t s02, const int64_t s03,
+        const int64_t s10, const int64_t s11, const int64_t s12,
+        const int innerq_is_k,
+        const int64_t s1,  const int64_t s2,  const int64_t s3,
+        const uint3 ne00_fd, const uint3 ne01_fd, const uint3 ne02_fd,
+        const uint3 ne11_fd, const uint3 ne12_fd) {
+
+    const int64_t group = blockIdx.x;
+    if (group >= ne_total_groups) return;
+
+    const int sid = threadIdx.x; // state index 0..511
+
+    // Compute source and destination pointers (same index math as turbo3)
+    const int64_t i_base = group * QK_TURBO3_TCQ;
+    uint32_t tmp = (uint32_t)i_base; uint2 div_mod;
+    div_mod = fast_div_modulo(tmp, ne00_fd); const int64_t i00 = div_mod.y; tmp = div_mod.x;
+    div_mod = fast_div_modulo(tmp, ne01_fd); const int64_t i01 = div_mod.y; tmp = div_mod.x;
+    div_mod = fast_div_modulo(tmp, ne02_fd); const int64_t i02 = div_mod.y; const int64_t i03 = div_mod.x;
+    const int64_t i12 = fastmodulo((uint32_t)i03, ne12_fd);
+    const int64_t i11 = fastmodulo((uint32_t)i02, ne11_fd);
+    const int64_t dst_row = *(src1 + i01*s10 + i11*s11 + i12*s12);
+    const float * grp_src = src0 + i01*s01 + i02*s02 + i03*s03 + i00;
+    block_turbo3_tcq * dst_blk = (block_turbo3_tcq *)((char *)dst + dst_row*s1 + i02*s2 + i03*s3)
+                                  + (i00 / QK_TURBO3_TCQ);
+
+    // Shared memory layout (~5KB, was ~35KB before global bt optimization):
+    // x[128]     : rotated+normalized input (also reused as outputs[] after Viterbi)
+    // cost[512]  : path costs buffer A (also reused for reductions)
+    // cost_b[512]: path costs buffer B (double-buffering eliminates 2/3 of syncs)
+    // Backtrace: one predecessor byte for each of the 64 low-state groups per
+    // step. The predecessor is independent of the output bits in sid[8:6], so
+    // storing 128×64 bytes is equivalent to the older 128×512 layout.
+    extern __shared__ uint8_t bt_shared[];
+    __shared__ acc_t x[128];
+    __shared__ acc_t cost[512];
+    __shared__ acc_t cost_b[512];
+    __shared__ int warp_min_idx[16];
+    __shared__ acc_t warp_min_cost[16];
+    __shared__ acc_t pred_min_cost[64];
+    __shared__ int shared_initial_state;
+
+    if (sid < 128) x[sid] = grp_src[sid];
+    __syncthreads();
+
+    __syncthreads();
+
+    // Norm reduction
+    cost[sid] = (sid < 128) ? x[sid] * x[sid] : 0.0f;
+    __syncthreads();
+    for (int stride = 256; stride >= 32; stride >>= 1) {
+        if (sid < stride) cost[sid] += cost[sid + stride];
+        __syncthreads();
+    }
+    if (sid < 32) {
+        acc_t v = cost[sid];
+        v += __shfl_down_sync(0xFFFFFFFFULL, v, 16);
+        v += __shfl_down_sync(0xFFFFFFFFULL, v, 8);
+        v += __shfl_down_sync(0xFFFFFFFFULL, v, 4);
+        v += __shfl_down_sync(0xFFFFFFFFULL, v, 2);
+        v += __shfl_down_sync(0xFFFFFFFFULL, v, 1);
+        if (sid == 0) cost[0] = v;
+    }
+    __syncthreads();
+    acc_t grp_norm = tcq_acc_sqrt(cost[0]);
+    acc_t inv_norm = grp_norm > 1e-10f ? 1.0f / grp_norm : 0.0f;
+
+    if (sid < 128) x[sid] *= inv_norm;
+    __syncthreads();
+
+    // FWHT. The first five stages are contained within each 32-lane warp, so
+    // use warp shuffles and only synchronize for the two cross-warp stages.
+    if (sid < 128) {
+        acc_t v = x[sid] * TURBO3_WHT_SIGNS1[sid];
+        const int lane = sid & 31;
+#pragma unroll
+        for (int h = 1; h < 32; h <<= 1) {
+            const acc_t other = __shfl_xor_sync(0xFFFFFFFFULL, v, h);
+            v = (lane & h) ? (other - v) : (v + other);
+        }
+        x[sid] = v;
+    }
+    __syncthreads();
+    if (sid < 64) {
+        const int j = ((sid >> 5) << 6) + (sid & 31);
+        acc_t a = x[j], b = x[j + 32];
+        x[j] = a + b; x[j + 32] = a - b;
+    }
+    __syncthreads();
+    if (sid < 64) {
+        acc_t a = x[sid], b = x[sid + 64];
+        x[sid] = a + b; x[sid + 64] = a - b;
+    }
+    __syncthreads();
+    constexpr float inv_sqrt_128 = 0.08838834764831845f;
+    if (sid < 128) x[sid] *= inv_sqrt_128 * TURBO3_WHT_SIGNS2[sid];
+    __syncthreads();
+
+    // opencoti-hook: TCQ activation-dump (F-D, #515). x[0..127] is now the post-FWHT unit-norm whitened vec —
+    // exactly the xt the trellis quantizes against d_turbo3_tcq_codebook below, and x stays read-only through the
+    // forward pass. Stage it for the F-E codebook retrain when enabled; nullptr (default) is a uniform no-op.
+    // group == blockIdx.x < ne_total_groups, so [group*128 + sid] is in-bounds of the (ne_total_groups*128) buffer.
+    if (d_tcq_dump_buf != nullptr && sid < 128) {
+        d_tcq_dump_buf[group * 128 + sid] = (float) x[sid];
+    }
+
+    if (sid == 0) cost[0] = grp_norm;
+    __syncthreads();
+
+    acc_t saved_norm = cost[0];
+
+    // Viterbi forward pass: double-buffered cost (1 sync/step, was 3)
+    uint8_t * bt = use_shared_bt ? bt_shared : bt_buf + (int64_t)blockIdx.x * (128 * 64);
+    cost[sid] = 0.0f;
+    __syncthreads();
+
+    for (int t = 0; t < 128; t++) {
+        // Double-buffer: even steps read cost/write cost_b, odd steps read cost_b/write cost
+        acc_t * cost_rd = (t & 1) ? cost_b : cost;
+        acc_t * cost_wr = (t & 1) ? cost   : cost_b;
+
+        acc_t xt = x[t];
+
+        // Right-shift trellis: ns = (prev >> 3) | (out << 6). The best
+        // predecessor depends only on sid's low 6 bits, so compute those 64
+        // minima once instead of repeating the same 8-way scan for each out.
+        if (sid < 64) {
+            const int base_prev = sid << 3;
+            acc_t best = cost_rd[base_prev];
+            int best_p = 0;
+#pragma unroll
+            for (int p = 1; p < 8; p++) {
+                acc_t c = cost_rd[base_prev | p];
+                if (c < best) {
+                    best = c;
+                    best_p = p;
+                }
+            }
+            pred_min_cost[sid] = best;
+            bt[t * 64 + sid] = (uint8_t) best_p;
+        }
+        __syncthreads();
+
+        const int pred_idx = sid & 0x3F;
+        acc_t dist = xt - d_turbo3_tcq_codebook[sid];
+        dist = dist * dist;
+
+        cost_wr[sid] = pred_min_cost[pred_idx] + dist;
+        __syncthreads();
+    }
+    // After 128 steps (even count): final costs are in cost[] (step 127 writes to cost)
+
+    // Warp argmin over 512 costs
+    {
+        acc_t my_cost = cost[sid];
+        int my_idx = sid;
+        #pragma unroll
+        for (int offset = 16; offset > 0; offset >>= 1) {
+            acc_t other_cost = __shfl_xor_sync(0xFFFFFFFFULL, my_cost, offset);
+            int other_idx = __shfl_xor_sync(0xFFFFFFFFULL, my_idx, offset);
+            if (other_cost < my_cost) { my_cost = other_cost; my_idx = other_idx; }
+        }
+        if (sid % 32 == 0) {
+            warp_min_cost[sid / 32] = my_cost;
+            warp_min_idx[sid / 32] = my_idx;
+        }
+    }
+    __syncthreads();
+    if (sid < 32) {
+        acc_t best = (sid < 16) ? warp_min_cost[sid] : 3.4028234663852886e38f;
+        int best_idx = (sid < 16) ? warp_min_idx[sid] : 0;
+#pragma unroll
+        for (int offset = 16; offset > 0; offset >>= 1) {
+            acc_t other_cost = __shfl_down_sync(0xFFFFFFFFULL, best, offset);
+            int other_idx = __shfl_down_sync(0xFFFFFFFFULL, best_idx, offset);
+            if (other_cost < best) {
+                best = other_cost;
+                best_idx = other_idx;
+            }
+        }
+        if (sid == 0) {
+            shared_initial_state = best_idx; // temporarily: best final state (becomes initial after backtrack)
+        }
+    }
+    __syncthreads();
+
+    // Save x[] to global buffer before backtrack overwrites it
+
+    // Backtrack (inherently sequential, reads global bt)
+    uint8_t * outputs = (uint8_t *)x;
+    if (sid == 0) {
+        int state = shared_initial_state;
+        for (int t = 127; t >= 0; t--) {
+            outputs[t] = (uint8_t)(state >> 6);
+            int p = bt[t * 64 + (state & 0x3F)];
+            state = ((state & 0x3F) << 3) | p;
+        }
+        shared_initial_state = state;
+    }
+    __syncthreads();
+
+    // Save output symbols to global buffer
+
+    // Parallel recon norm: t>=2 can compute state directly from 3 outputs (3 shifts of 3 = 9 bits)
+    acc_t my_recon_sq = 0.0f;
+    if (sid < 128) {
+        int cur_state;
+        if (sid < 2) {
+            cur_state = shared_initial_state;
+            for (int t = 0; t <= sid; t++)
+                cur_state = (cur_state >> 3) | (((int)outputs[t]) << 6);
+        } else {
+            cur_state = ((int)outputs[sid - 2] & 0x7)
+                      | (((int)outputs[sid - 1] & 0x7) << 3)
+                      | (((int)outputs[sid]     & 0x7) << 6);
+        }
+        acc_t c = d_turbo3_tcq_codebook[cur_state];
+        my_recon_sq = c * c;
+    }
+    cost[sid] = my_recon_sq;
+    __syncthreads();
+    for (int stride = 256; stride >= 32; stride >>= 1) {
+        if (sid < stride) cost[sid] += cost[sid + stride];
+        __syncthreads();
+    }
+    if (sid < 32) {
+        acc_t v = cost[sid];
+        v += __shfl_down_sync(0xFFFFFFFFULL, v, 16);
+        v += __shfl_down_sync(0xFFFFFFFFULL, v, 8);
+        v += __shfl_down_sync(0xFFFFFFFFULL, v, 4);
+        v += __shfl_down_sync(0xFFFFFFFFULL, v, 2);
+        v += __shfl_down_sync(0xFFFFFFFFULL, v, 1);
+        if (sid == 0) cost[0] = v;
+    }
+    __syncthreads();
+    acc_t recon_norm = tcq_acc_sqrt(cost[0]);
+    acc_t corrected_norm = (recon_norm > 1e-10f) ? saved_norm / recon_norm : saved_norm;
+    corrected_norm *= innerq_is_k ? d_tcq_norm_alpha : d_tcq_norm_alpha_v;
+
+    // Parallel bitpack: qs stores 6 initial-state bits followed by 128 3-bit
+    // output symbols. Each byte is independent, so avoid the old serial OR loop.
+    if (sid < 49) {
+        const int init_bits = (shared_initial_state >> 3) & 0x3F;
+        uint8_t packed = 0;
+#pragma unroll
+        for (int bit = 0; bit < 8; bit++) {
+            const int pos = sid * 8 + bit;
+            int v = 0;
+            if (pos < 6) {
+                v = (init_bits >> pos) & 1;
+            } else {
+                const int sym_bit_pos = pos - 6;
+                const int sym_idx = sym_bit_pos / 3;
+                if (sym_idx < 128) {
+                    v = (outputs[sym_idx] >> (sym_bit_pos % 3)) & 1;
+                }
+            }
+            packed |= (uint8_t)(v << bit);
+        }
+        dst_blk->qs[sid] = packed;
+    }
+    if (sid == 0) {
+        dst_blk->norm = __float2half((float)corrected_norm);
+    }
+}
+
+// TCQ GET_ROWS dequantize (for non-FA paths)
+#define QR_TURBO3_TCQ 2
+static __device__ __forceinline__
+void dequantize_turbo3_tcq(const void * vx, const int64_t ib, const int iqs, float2 & v) {
+    const block_turbo3_tcq * blk = (const block_turbo3_tcq *)vx + ib;
+    const float norm = __half2float(blk->norm);
+
+    // Decode element iqs
+    {
+        const int t = iqs;
+        const int bit_pos = t * 3;
+        const int byte_idx = bit_pos / 8;
+        const int bit_off = bit_pos % 8;
+        const uint16_t raw = (uint16_t)blk->qs[byte_idx] | ((uint16_t)blk->qs[byte_idx + 1] << 8);
+        const int state = (raw >> bit_off) & 0x1FF;
+        v.x = d_turbo3_tcq_codebook[state] * norm;
+    }
+    // Decode element iqs + 64 (stride = half block size)
+    {
+        const int t = iqs + 64;
+        const int bit_pos = t * 3;
+        const int byte_idx = bit_pos / 8;
+        const int bit_off = bit_pos % 8;
+        const uint16_t raw = (uint16_t)blk->qs[byte_idx] | ((uint16_t)blk->qs[byte_idx + 1] << 8);
+        const int state = (raw >> bit_off) & 0x1FF;
+        v.y = d_turbo3_tcq_codebook[state] * norm;
+    }
+}
+
+// =====================================================================================
+// TURBO2_TCQ: 2-bit Trellis-Coded Quantization (k=2, L=8, 256 states, free initial state)
+// =====================================================================================
+
+// 2-bit TCQ codebook (product_mono/iter090, 256-state bitshift trellis). If you copy these, credit spiritbuun!
+// CUDA GLA product-aware training, 100 iters on Qwen3.5-27B FWHT-rotated KV activations. Decode: state_t = read_8_bits(qs, t*2)
+static __constant__ float d_turbo2_tcq_codebook[256] = {
+    -0.18030643f, -0.11009848f, -0.04742626f, +0.02894132f, -0.10523465f, -0.031312924f, +0.031491395f, +0.12263535f,
+    -0.15660362f, -0.055477407f, +0.0046675834f, +0.06166081f, -0.07506216f, -0.016963918f, +0.043737844f, +0.116496615f,
+    -0.08632783f, -0.022493735f, +0.041032985f, +0.10660284f, -0.06274858f, -0.0036939639f, +0.02095157f, +0.07539709f,
+    -0.09802641f, -0.008419088f, +0.059072323f, +0.17311879f, -0.093109086f, -0.02654333f, +0.014827672f, +0.07793592f,
+    -0.031235758f, +0.01271591f, +0.08752262f, +0.17246453f, -0.14595252f, -0.07227624f, +0.013628688f, +0.08131674f,
+    -0.036909282f, +0.0018896917f, +0.05209119f, +0.12407892f, -0.13689458f, -0.06054520f, +0.0064648795f, +0.07551241f,
+    -0.18980840f, -0.110128626f, -0.046503957f, +0.026387159f, -0.034967307f, +0.04810357f, +0.072072044f, +0.14355458f,
+    -0.10182410f, -0.02907887f, +0.014033012f, +0.083419636f, -0.056140676f, +0.008405868f, +0.066070884f, +0.14037225f,
+    -0.117427245f, -0.047159385f, +0.016928354f, +0.08142885f, -0.029359628f, +0.045608785f, +0.10559447f, +0.20061271f,
+    -0.040425077f, +0.029068163f, +0.08408973f, +0.13628258f, -0.16633821f, -0.10711727f, -0.04196669f, +0.027895834f,
+    -0.0054065837f, +0.058898676f, +0.12688550f, +0.18268861f, -0.16287325f, -0.11218357f, -0.07165227f, -0.009524379f,
+    -0.24026902f, -0.073219374f, -0.0005165726f, +0.05959821f, -0.05532953f, +0.027044486f, +0.09425678f, +0.15356481f,
+    -0.14381111f, -0.10563502f, -0.037867088f, +0.023611993f, -0.03624307f, +0.049588434f, +0.12192037f, +0.23462485f,
+    -0.14990251f, -0.09659304f, -0.05886742f, +0.014878461f, -0.009889551f, +0.06910514f, +0.12120181f, +0.22596690f,
+    -0.08290075f, -0.009009629f, +0.066151775f, +0.12188313f, -0.11591514f, -0.06952189f, -0.031633306f, +0.023740824f,
+    -0.20510401f, -0.103369795f, +0.09148037f, +0.17268716f, -0.16597997f, -0.09207068f, -0.032810967f, +0.024847647f,
+    -0.02487482f, +0.049298953f, +0.09624215f, +0.14217524f, -0.18418685f, -0.10147012f, -0.05841265f, +0.008057022f,
+    -0.14269894f, -0.092456274f, -0.026881337f, +0.049792137f, -0.019881032f, +0.030333601f, +0.09736802f, +0.17764080f,
+    -0.19579841f, -0.114739306f, -0.026823774f, +0.07466014f, -0.09001050f, -0.041468445f, +0.028473806f, +0.08870695f,
+    -0.019396419f, +0.042828932f, +0.10885327f, +0.13335012f, -0.15005013f, -0.074581385f, -0.028608415f, +0.03848942f,
+    -0.09687270f, -0.057059396f, +0.0077843578f, +0.06302297f, -0.23247094f, -0.14509225f, -0.032651436f, +0.027010715f,
+    -0.047595482f, +0.06280303f, +0.114691675f, +0.17124057f, -0.21092793f, -0.13704823f, -0.07340412f, +0.0039013291f,
+    -0.062834196f, +0.012601906f, +0.012601906f, +0.08721347f, -0.13256435f, -0.024173854f, +0.07723171f, +0.14801070f,
+    -0.06471605f, -0.0017903054f, -0.0017903054f, +0.058302354f, -0.09731802f, -0.03400696f, +0.02762442f, +0.08986137f,
+    -0.08288722f, -0.019051429f, +0.045709886f, +0.15211061f, -0.09507891f, -0.015612489f, +0.025347246f, +0.087257534f,
+    -0.066236064f, -0.0047936034f, +0.06386274f, +0.15401669f, -0.105809286f, -0.051802177f, +0.01073050f, +0.08292137f,
+    -0.11884470f, -0.04404144f, +0.02550729f, +0.02550729f, -0.01731189f, +0.062161792f, +0.12127554f, +0.21981733f,
+    -0.17066145f, -0.11660990f, -0.049425896f, +0.021293938f, -0.04711412f, +0.026577346f, +0.055197213f, +0.12541275f,
+    -0.028268812f, +0.015206398f, +0.09002519f, +0.12699963f, -0.10059831f, -0.026676945f, +0.059903253f, +0.13054545f,
+    -0.09582803f, -0.033371232f, +0.010346129f, +0.066766635f, -0.09964944f, -0.028686784f, +0.021184925f, +0.09120017f,
+    -0.16957201f, -0.07594450f, +0.04172865f, +0.18313301f, -0.051526368f, +0.011877304f, +0.011877304f, +0.07956263f,
+    -0.13432936f, -0.05269006f, +0.03536416f, +0.117640756f, -0.022776067f, +0.042032316f, +0.10472976f, +0.18042557f
+};
+
+// 2-bit TCQ SET_ROWS encode: Viterbi optimal path with right-shift trellis (k=2, L=8)
+// Double-buffered cost arrays + global memory backtrace (128 syncs/group, was 384)
+template<typename idx_t, typename acc_t = float>
+static __global__ void __launch_bounds__(256, 1) k_set_rows_turbo2_tcq(
+        const float * __restrict__ src0, const idx_t * __restrict__ src1,
+        block_turbo2_tcq * __restrict__ dst, const int64_t ne_total_groups,
+        uint8_t * __restrict__ bt_buf,
+        const int use_shared_bt,
+        const int64_t ne00, const int64_t ne01, const int64_t ne02,
+        const int64_t ne10, const int64_t ne11, const int64_t ne12, const int64_t ne13,
+        const int64_t s01, const int64_t s02, const int64_t s03,
+        const int64_t s10, const int64_t s11, const int64_t s12,
+        const int iq_is_k,
+        const int64_t s1, const int64_t s2, const int64_t s3,
+        const uint3 ne00_fd, const uint3 ne01_fd, const uint3 ne02_fd,
+        const uint3 ne11_fd, const uint3 ne12_fd) {
+
+    const int grp = blockIdx.x;
+    if (grp >= ne_total_groups) return;
+    const int sid = threadIdx.x; // 0..255 = trellis state
+
+    // Compute source and destination pointers (all threads, used by thread 0)
+    const int64_t i_base = int64_t(grp) * QK_TURBO2_TCQ;
+    uint32_t tmp = (uint32_t)i_base; uint2 div_mod;
+    div_mod = fast_div_modulo(tmp, ne00_fd); const int64_t i00 = div_mod.y; tmp = div_mod.x;
+    div_mod = fast_div_modulo(tmp, ne01_fd); const int64_t i01 = div_mod.y; tmp = div_mod.x;
+    div_mod = fast_div_modulo(tmp, ne02_fd); const int64_t i02 = div_mod.y; const int64_t i03 = div_mod.x;
+    const int64_t i12 = fastmodulo((uint32_t)i03, ne12_fd);
+    const int64_t i11 = fastmodulo((uint32_t)i02, ne11_fd);
+    const int64_t dst_row = *(src1 + i01*s10 + i11*s11 + i12*s12);
+    const float * grp_src = src0 + i01*s01 + i02*s02 + i03*s03 + i00;
+    block_turbo2_tcq * dst_blk = (block_turbo2_tcq *)((char *)dst + dst_row*s1 + i02*s2 + i03*s3)
+                               + (i00 / QK_TURBO2_TCQ);
+
+    // Backtrace: one predecessor byte per 64 low-state groups per step.
+    // The predecessor depends only on sid's low 6 bits (same as turbo3_tcq).
+    extern __shared__ uint8_t bt_shared[];
+    __shared__ acc_t x[128];
+    __shared__ acc_t cost[256];
+    __shared__ acc_t cost_b[256];   // double-buffering for Viterbi
+    __shared__ int warp_min_idx[8];
+    __shared__ acc_t warp_min_cost[8];
+    __shared__ acc_t pred_min_cost[64];
+    __shared__ int shared_initial_state;
+
+    if (sid < 128) x[sid] = grp_src[sid];
+    __syncthreads();
+
+    __syncthreads();
+
+    // Norm reduction
+    cost[sid] = (sid < 128) ? x[sid] * x[sid] : 0.0f;
+    __syncthreads();
+    for (int stride = 128; stride >= 32; stride >>= 1) {
+        if (sid < stride) cost[sid] += cost[sid + stride];
+        __syncthreads();
+    }
+    if (sid < 32) {
+        acc_t v = cost[sid];
+        v += __shfl_down_sync(0xFFFFFFFFULL, v, 16);
+        v += __shfl_down_sync(0xFFFFFFFFULL, v, 8);
+        v += __shfl_down_sync(0xFFFFFFFFULL, v, 4);
+        v += __shfl_down_sync(0xFFFFFFFFULL, v, 2);
+        v += __shfl_down_sync(0xFFFFFFFFULL, v, 1);
+        if (sid == 0) cost[0] = v;
+    }
+    __syncthreads();
+    acc_t grp_norm = tcq_acc_sqrt(cost[0]);
+    acc_t inv_norm = grp_norm > 1e-10f ? 1.0f / grp_norm : 0.0f;
+
+    if (sid < 128) x[sid] *= inv_norm;
+    __syncthreads();
+
+    // FWHT. The first five stages use warp shuffles, the two cross-warp stages
+    // use shared memory (same approach as turbo3_tcq).
+    if (sid < 128) {
+        acc_t v = x[sid] * TURBO3_WHT_SIGNS1[sid];
+        const int lane = sid & 31;
+#pragma unroll
+        for (int h = 1; h < 32; h <<= 1) {
+            const acc_t other = __shfl_xor_sync(0xFFFFFFFFULL, v, h);
+            v = (lane & h) ? (other - v) : (v + other);
+        }
+        x[sid] = v;
+    }
+    __syncthreads();
+    if (sid < 64) {
+        const int j = ((sid >> 5) << 6) + (sid & 31);
+        acc_t a = x[j], b = x[j + 32];
+        x[j] = a + b; x[j + 32] = a - b;
+    }
+    __syncthreads();
+    if (sid < 64) {
+        acc_t a = x[sid], b = x[sid + 64];
+        x[sid] = a + b; x[sid + 64] = a - b;
+    }
+    __syncthreads();
+    constexpr float inv_sqrt_128 = 0.08838834764831845f;
+    if (sid < 128) x[sid] *= inv_sqrt_128 * TURBO3_WHT_SIGNS2[sid];
+    __syncthreads();
+
+    if (sid == 0) cost[0] = grp_norm;
+    __syncthreads();
+
+    acc_t saved_norm = cost[0];
+
+    // Viterbi forward pass: double-buffered cost (1 sync/step, was 3)
+    uint8_t * bt = use_shared_bt ? bt_shared : bt_buf + (int64_t)blockIdx.x * (128 * 64);
+    cost[sid] = 0.0f;
+    __syncthreads();
+
+    for (int t = 0; t < 128; t++) {
+        acc_t * cost_rd = (t & 1) ? cost_b : cost;
+        acc_t * cost_wr = (t & 1) ? cost   : cost_b;
+
+        acc_t xt = x[t];
+
+        // Right-shift trellis (k=2, L=8): ns = (prev >> 2) | (out << 6).
+        // The best predecessor depends only on sid's low 6 bits, so compute
+        // those 64 minima once instead of repeating the same 4-way scan per state.
+        if (sid < 64) {
+            const int base_prev = sid << 2;
+            acc_t best = cost_rd[base_prev];
+            int best_p = 0;
+#pragma unroll
+            for (int p = 1; p < 4; p++) {
+                acc_t c = cost_rd[base_prev | p];
+                if (c < best) {
+                    best = c;
+                    best_p = p;
+                }
+            }
+            pred_min_cost[sid] = best;
+            bt[t * 64 + sid] = (uint8_t) best_p;
+        }
+        __syncthreads();
+
+        const int pred_idx = sid & 0x3F;
+        acc_t dist = xt - d_turbo2_tcq_codebook[sid];
+        dist = dist * dist;
+
+        cost_wr[sid] = pred_min_cost[pred_idx] + dist;
+        __syncthreads();
+    }
+    // After 128 steps (even count): final costs in cost[]
+
+    // Warp argmin over 256 costs
+    {
+        acc_t my_cost = cost[sid];
+        int my_idx = sid;
+        #pragma unroll
+        for (int offset = 16; offset > 0; offset >>= 1) {
+            acc_t other_cost = __shfl_xor_sync(0xFFFFFFFFULL, my_cost, offset);
+            int other_idx = __shfl_xor_sync(0xFFFFFFFFULL, my_idx, offset);
+            if (other_cost < my_cost) { my_cost = other_cost; my_idx = other_idx; }
+        }
+        if (sid % 32 == 0) {
+            warp_min_cost[sid / 32] = my_cost;
+            warp_min_idx[sid / 32] = my_idx;
+        }
+    }
+    __syncthreads();
+    if (sid < 32) {
+        acc_t best = (sid < 8) ? warp_min_cost[sid] : 3.4028234663852886e38f;
+        int best_idx = (sid < 8) ? warp_min_idx[sid] : 0;
+#pragma unroll
+        for (int offset = 16; offset > 0; offset >>= 1) {
+            acc_t other_cost = __shfl_down_sync(0xFFFFFFFFULL, best, offset);
+            int other_idx = __shfl_down_sync(0xFFFFFFFFULL, best_idx, offset);
+            if (other_cost < best) {
+                best = other_cost;
+                best_idx = other_idx;
+            }
+        }
+        if (sid == 0) {
+            shared_initial_state = best_idx;
+        }
+    }
+    __syncthreads();
+
+    // Save x[] to global buffer before backtrack overwrites it
+
+    // Backtrack (inherently sequential, reads compressed bt)
+    uint8_t * outputs = (uint8_t *)x;
+    if (sid == 0) {
+        int state = shared_initial_state;
+        for (int t = 127; t >= 0; t--) {
+            outputs[t] = (uint8_t)(state >> 6);
+            int p = bt[t * 64 + (state & 0x3F)];
+            state = ((state & 0x3F) << 2) | p;
+        }
+        shared_initial_state = state;
+    }
+    __syncthreads();
+
+    // Save output symbols to global buffer
+
+    // Parallel recon norm: t>=3 can compute state directly from 4 outputs (4 shifts of 2 = 8 bits)
+    acc_t my_recon_sq = 0.0f;
+    if (sid < 128) {
+        int cur_state;
+        if (sid < 3) {
+            cur_state = shared_initial_state;
+            for (int t = 0; t <= sid; t++)
+                cur_state = (cur_state >> 2) | (((int)outputs[t]) << 6);
+        } else {
+            cur_state = ((int)outputs[sid - 3] & 0x3)
+                      | (((int)outputs[sid - 2] & 0x3) << 2)
+                      | (((int)outputs[sid - 1] & 0x3) << 4)
+                      | (((int)outputs[sid]     & 0x3) << 6);
+        }
+        acc_t c = d_turbo2_tcq_codebook[cur_state];
+        my_recon_sq = c * c;
+    }
+    cost[sid] = my_recon_sq;
+    __syncthreads();
+    for (int stride = 128; stride >= 32; stride >>= 1) {
+        if (sid < stride) cost[sid] += cost[sid + stride];
+        __syncthreads();
+    }
+    if (sid < 32) {
+        acc_t v = cost[sid];
+        v += __shfl_down_sync(0xFFFFFFFFULL, v, 16);
+        v += __shfl_down_sync(0xFFFFFFFFULL, v, 8);
+        v += __shfl_down_sync(0xFFFFFFFFULL, v, 4);
+        v += __shfl_down_sync(0xFFFFFFFFULL, v, 2);
+        v += __shfl_down_sync(0xFFFFFFFFULL, v, 1);
+        if (sid == 0) cost[0] = v;
+    }
+    __syncthreads();
+    acc_t recon_norm = tcq_acc_sqrt(cost[0]);
+    acc_t corrected_norm = (recon_norm > 1e-10f) ? saved_norm / recon_norm : saved_norm;
+    corrected_norm *= iq_is_k ? d_tcq_norm_alpha : d_tcq_norm_alpha_v;
+
+    // Parallel bitpack: qs stores 6 initial-state bits followed by 128 2-bit
+    // output symbols. Each byte is independent (2-bit symbols never cross byte
+    // boundaries after the 6-bit prefix), so avoid the serial OR loop.
+    if (sid < 33) {
+        const int init_bits = (shared_initial_state >> 2) & 0x3F;
+        uint8_t packed = 0;
+#pragma unroll
+        for (int bit = 0; bit < 8; bit++) {
+            const int pos = sid * 8 + bit;
+            int v = 0;
+            if (pos < 6) {
+                v = (init_bits >> pos) & 1;
+            } else {
+                const int sym_bit_pos = pos - 6;
+                const int sym_idx = sym_bit_pos / 2;
+                if (sym_idx < 128) {
+                    v = (outputs[sym_idx] >> (sym_bit_pos % 2)) & 1;
+                }
+            }
+            packed |= (uint8_t)(v << bit);
+        }
+        dst_blk->qs[sid] = packed;
+    }
+    if (sid == 0) {
+        dst_blk->norm = __float2half((float)corrected_norm);
+    }
+}
+
+// 2-bit TCQ GET_ROWS dequantize
+#define QR_TURBO2_TCQ 2
+static __device__ __forceinline__
+void dequantize_turbo2_tcq(const void * vx, const int64_t ib, const int iqs, float2 & v) {
+    const block_turbo2_tcq * blk = (const block_turbo2_tcq *)vx + ib;
+    const float norm = __half2float(blk->norm);
+
+    // Decode element iqs: read 8-bit state via sliding window
+    {
+        const int t = iqs;
+        const int bit_pos = t * 2;
+        const int byte_idx = bit_pos / 8;
+        const int bit_off = bit_pos % 8;
+        const uint16_t raw = (uint16_t)blk->qs[byte_idx] | ((uint16_t)blk->qs[byte_idx + 1] << 8);
+        const int state = (raw >> bit_off) & 0xFF;
+        v.x = d_turbo2_tcq_codebook[state] * norm;
+    }
+    // Decode element iqs + 64
+    {
+        const int t = iqs + 64;
+        const int bit_pos = t * 2;
+        const int byte_idx = bit_pos / 8;
+        const int bit_off = bit_pos % 8;
+        const uint16_t raw = (uint16_t)blk->qs[byte_idx] | ((uint16_t)blk->qs[byte_idx + 1] << 8);
+        const int state = (raw >> bit_off) & 0xFF;
+        v.y = d_turbo2_tcq_codebook[state] * norm;
+    }
+}
diff --git a/llama.cpp/ggml/src/ggml-quants.c b/llama.cpp/ggml/src/ggml-quants.c
index 15d231f70..3ac0f340a 100644
--- a/llama.cpp/ggml/src/ggml-quants.c
+++ b/llama.cpp/ggml/src/ggml-quants.c
@@ -190,6 +190,57 @@ void quantize_row_q5_0_ref(const float * GGML_RESTRICT x, block_q5_0 * GGML_REST
     }
 }
 
+// opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port) — see docs/features/poly_kv.md.
+// Signed 6-bit (levels offset by 32). qs[16] = low 4 bits of all 32 values;
+// qh[8] = high 2 bits, packed 4 values/byte. Ported from ik_llama ggml-quants.c:853.
+void quantize_row_q6_0_ref(const float * GGML_RESTRICT x, block_q6_0 * GGML_RESTRICT y, int64_t k) {
+    static const int qk = QK6_0;
+
+    assert(k % qk == 0);
+
+    const int nb = k / qk;
+
+    for (int i = 0; i < nb; i++) {
+        float amax = 0.0f; // absolute max
+        float max  = 0.0f;
+
+        for (int j = 0; j < qk; j++) {
+            const float v = x[i*qk + j];
+            if (amax < fabsf(v)) {
+                amax = fabsf(v);
+                max  = v;
+            }
+        }
+
+        const float d  = max / -32;
+        const float id = d ? 1.0f/d : 0.0f;
+
+        memset(y[i].qh, 0, qk/4);
+
+        float sumqx = 0, sumq2 = 0;
+        for (int j = 0; j < qk/2; ++j) {
+            const float x0 = x[i*qk + 0    + j]*id;
+            const float x1 = x[i*qk + qk/2 + j]*id;
+            const float w0 = x0*x0;
+            const float w1 = x1*x1;
+
+            const uint8_t xi0 = MIN(63, (int8_t)(x0 + 32.5f));
+            const uint8_t xi1 = MIN(63, (int8_t)(x1 + 32.5f));
+
+            y[i].qs[j] = (xi0 & 0x0F) | ((xi1 & 0x0F) << 4);
+
+            const uint8_t h = (xi0 >> 4) | ((xi1 >> 4) << 2);
+            y[i].qh[j%(qk/4)] |= (h << 4*(j/(qk/4)));
+
+            const float q0 = (float)xi0 - 32.f;
+            const float q1 = (float)xi1 - 32.f;
+            sumqx += w0*x[i*qk + j]*q0 + w1*x[i*qk + qk/2 + j]*q1;
+            sumq2 += w0*q0*q0 + w1*q1*q1;
+        }
+        y[i].d = sumq2 > 0 ? GGML_FP32_TO_FP16(sumqx/sumq2) : GGML_FP32_TO_FP16(d);
+    }
+}
+
 void quantize_row_q5_1_ref(const float * GGML_RESTRICT x, block_q5_1 * GGML_RESTRICT y, int64_t k) {
     const int qk = QK5_1;
 
@@ -465,6 +516,29 @@ void dequantize_row_q5_0(const block_q5_0 * GGML_RESTRICT x, float * GGML_RESTRI
     }
 }
 
+// opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port) — ported from ik_llama ggml-quants.c:1675.
+void dequantize_row_q6_0(const block_q6_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) {
+    static const int qk = QK6_0;
+
+    assert(k % qk == 0);
+
+    const int nb = k / qk;
+
+    for (int i = 0; i < nb; i++) {
+        const float d = GGML_FP16_TO_FP32(x[i].d);
+
+        for (int j = 0; j < qk/2; ++j) {
+            const uint8_t h = x[i].qh[j%(qk/4)] >> 4*(j/(qk/4));
+
+            const int32_t x0 = ((x[i].qs[j] & 0x0F) | ((h << 4) & 0x30)) - 32;
+            const int32_t x1 = ((x[i].qs[j] >>   4) | ((h << 2) & 0x30)) - 32;
+
+            y[i*qk + j + 0   ] = x0*d;
+            y[i*qk + j + qk/2] = x1*d;
+        }
+    }
+}
+
 void dequantize_row_q5_1(const block_q5_1 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) {
     static const int qk = QK5_1;
 
@@ -2167,6 +2241,58 @@ size_t quantize_q5_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst,
     return nrow * row_size;
 }
 
+// opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port) — KV path (quant_weights==NULL) routes to the
+// byte-exact ref; imatrix path mirrors ik_llama ggml-quants.c:3613 for GGUF-weight completeness.
+static void quantize_row_q6_0_impl(const float * GGML_RESTRICT x, block_q6_0 * GGML_RESTRICT y, int64_t n_per_row, const float * quant_weights) {
+    static_assert(QK6_0 == 32, "QK6_0 must be 32");
+
+    if (!quant_weights) {
+        quantize_row_q6_0_ref(x, y, n_per_row);
+        return;
+    }
+
+    float weight[QK6_0];
+    int8_t L[QK6_0];
+
+    float sum_x2 = 0;
+    for (int j = 0; j < n_per_row; ++j) sum_x2 += x[j]*x[j];
+    float sigma2 = sum_x2/n_per_row;
+
+    const int64_t nb = n_per_row/QK6_0;
+    for (int ib = 0; ib < nb; ++ib) {
+        const float * xb = x + QK6_0 * ib;
+        const float * qw = quant_weights + QK6_0 * ib;
+        for (int j = 0; j < QK6_0; ++j) weight[j] = qw[j] * sqrtf(sigma2 + xb[j]*xb[j]);
+        float d = make_qx_quants(QK6_0, 32, xb, L, 1, weight);
+        y[ib].d = GGML_FP32_TO_FP16(d);
+
+        memset(y[ib].qh, 0, QK6_0/4);
+
+        for (int j = 0; j < 16; ++j) {
+            const uint8_t xi0 = L[j];
+            const uint8_t xi1 = L[j+16];
+            y[ib].qs[j] = (xi0 & 0x0F) | ((xi1 & 0x0F) << 4);
+            const uint8_t h = (xi0 >> 4) | ((xi1 >> 4) << 2);
+            y[ib].qh[j%8] |= (h << 4*(j/8));
+        }
+    }
+}
+
+size_t quantize_q6_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) {
+    if (!quant_weights) {
+        quantize_row_q6_0_ref(src, dst, (int64_t)nrow*n_per_row);
+        return nrow * ggml_row_size(GGML_TYPE_Q6_0, n_per_row);
+    }
+    size_t row_size = ggml_row_size(GGML_TYPE_Q6_0, n_per_row);
+    char * qrow = (char *)dst;
+    for (int64_t row = 0; row < nrow; ++row) {
+        quantize_row_q6_0_impl(src, (block_q6_0*)qrow, n_per_row, quant_weights);
+        src += n_per_row;
+        qrow += row_size;
+    }
+    return nrow * row_size;
+}
+
 static void quantize_row_q5_1_impl(const float * GGML_RESTRICT x, block_q5_1 * GGML_RESTRICT y, int64_t n_per_row, const float * quant_weights) {
     static_assert(QK5_1 == 32, "QK5_1 must be 32");
 
diff --git a/llama.cpp/ggml/src/ggml-quants.h b/llama.cpp/ggml/src/ggml-quants.h
index 75c6efa82..d77693e92 100644
--- a/llama.cpp/ggml/src/ggml-quants.h
+++ b/llama.cpp/ggml/src/ggml-quants.h
@@ -21,6 +21,8 @@ GGML_API void quantize_row_q5_0_ref(const float * GGML_RESTRICT x, block_q5_0 *
 GGML_API void quantize_row_q5_1_ref(const float * GGML_RESTRICT x, block_q5_1 * GGML_RESTRICT y, int64_t k);
 GGML_API void quantize_row_q8_0_ref(const float * GGML_RESTRICT x, block_q8_0 * GGML_RESTRICT y, int64_t k);
 GGML_API void quantize_row_q8_1_ref(const float * GGML_RESTRICT x, block_q8_1 * GGML_RESTRICT y, int64_t k);
+// opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port)
+GGML_API void quantize_row_q6_0_ref(const float * GGML_RESTRICT x, block_q6_0 * GGML_RESTRICT y, int64_t k);
 
 GGML_API void quantize_row_mxfp4_ref(const float * GGML_RESTRICT x, block_mxfp4 * GGML_RESTRICT y, int64_t k);
 GGML_API void quantize_row_nvfp4_ref(const float * GGML_RESTRICT x, block_nvfp4 * GGML_RESTRICT y, int64_t k);
@@ -49,6 +51,8 @@ GGML_API void dequantize_row_q5_0(const block_q5_0 * GGML_RESTRICT x, float * GG
 GGML_API void dequantize_row_q5_1(const block_q5_1 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
 GGML_API void dequantize_row_q8_0(const block_q8_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
 //GGML_API void dequantize_row_q8_1(const block_q8_1 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
+// opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port)
+GGML_API void dequantize_row_q6_0(const block_q6_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
 
 GGML_API void dequantize_row_mxfp4(const block_mxfp4 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
 GGML_API void dequantize_row_nvfp4(const block_nvfp4 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
@@ -98,6 +102,8 @@ GGML_API size_t quantize_q4_1(const float * GGML_RESTRICT src, void * GGML_RESTR
 GGML_API size_t quantize_q5_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
 GGML_API size_t quantize_q5_1(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
 GGML_API size_t quantize_q8_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
+// opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port)
+GGML_API size_t quantize_q6_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
 
 GGML_API size_t quantize_mxfp4(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
 GGML_API size_t quantize_nvfp4(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
@@ -121,6 +127,14 @@ GGML_API size_t quantize_turbo2_0(const float * GGML_RESTRICT src, void * GGML_R
 GGML_API void quantize_row_turbo8_0_ref(const float * GGML_RESTRICT x, block_turbo8_0 * GGML_RESTRICT y, int64_t k);
 GGML_API void dequantize_row_turbo8_0(const block_turbo8_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
 GGML_API size_t quantize_turbo8_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
+// opencoti-hook: TCQ (F5 M6 C5) — Trellis-Coded Quantization KV (buun-aligned names). CPU stubs only;
+// real encode/decode live in CUDA (set-rows / FA-VEC). See ggml-turbo-quant.c.
+GGML_API void   quantize_row_turbo3_tcq_ref(const float * GGML_RESTRICT x, block_turbo3_tcq * GGML_RESTRICT y, int64_t k);
+GGML_API void   dequantize_row_turbo3_tcq(const block_turbo3_tcq * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
+GGML_API size_t quantize_turbo3_tcq(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
+GGML_API void   quantize_row_turbo2_tcq_ref(const float * GGML_RESTRICT x, block_turbo2_tcq * GGML_RESTRICT y, int64_t k);
+GGML_API void   dequantize_row_turbo2_tcq(const block_turbo2_tcq * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
+GGML_API size_t quantize_turbo2_tcq(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
 
 // TQ3_1S: WHT-rotated 3-bit weight quantization (8-level Lloyd-Max)
 GGML_API void quantize_row_tq3_1s_ref(const float * GGML_RESTRICT x, block_tq3_1s * GGML_RESTRICT y, int64_t k);
diff --git a/llama.cpp/ggml/src/ggml-turbo-quant.c b/llama.cpp/ggml/src/ggml-turbo-quant.c
index 38e15feca..156502191 100644
--- a/llama.cpp/ggml/src/ggml-turbo-quant.c
+++ b/llama.cpp/ggml/src/ggml-turbo-quant.c
@@ -652,6 +652,91 @@ size_t quantize_turbo8_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT d
     return nrows * row_size;
 }
 
+/* ---------- TURBO3_TCQ / TURBO2_TCQ: Trellis-Coded Quantization (CPU stubs) ----------
+ * opencoti-hook: TCQ (F5 M6 C5), ported verbatim from buun (spiritbuun/buun-llama-cpp).
+ * TCQ is GPU-only: the real encode (FWHT + Viterbi) is the CUDA SET_ROWS kernel and the real
+ * decode (O(1) sliding-window codebook lookup) is inlined in the FA-VEC kernel. These CPU paths
+ * are deliberate stubs — from_float writes only the corrected group L2 norm (so the
+ * dequant-on-lift / POSITION_WINDOW host-pinned tail keeps a consistent norm), to_float zero-fills.
+ */
+void quantize_row_turbo3_tcq_ref(const float * GGML_RESTRICT x, block_turbo3_tcq * GGML_RESTRICT y, int64_t k) {
+    // Stub — CUDA kernel handles TCQ quantize (Viterbi). CPU path zeros out.
+    assert(k % QK_TURBO3_TCQ == 0);
+    const int nb = k / QK_TURBO3_TCQ;
+    for (int i = 0; i < nb; i++) {
+        float norm = 0.0f;
+        for (int j = 0; j < QK_TURBO3_TCQ; j++) norm += x[i*QK_TURBO3_TCQ + j] * x[i*QK_TURBO3_TCQ + j];
+        y[i].norm = GGML_FP32_TO_FP16(sqrtf(norm));
+        memset(y[i].qs, 0, 49);
+    }
+}
+
+void dequantize_row_turbo3_tcq(const block_turbo3_tcq * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) {
+    GGML_UNUSED(x);
+    assert(k % QK_TURBO3_TCQ == 0);
+    const int nb = k / QK_TURBO3_TCQ;
+    for (int block = 0; block < nb; block++) {
+        for (int j = 0; j < QK_TURBO3_TCQ; j++) {
+            y[block * QK_TURBO3_TCQ + j] = 0.0f;
+        }
+    }
+}
+
+size_t quantize_turbo3_tcq(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst,
+                         int64_t nrows, int64_t n_per_row, const float * imatrix) {
+    GGML_UNUSED(imatrix);
+    assert(n_per_row % QK_TURBO3_TCQ == 0);
+
+    size_t row_size = (n_per_row / QK_TURBO3_TCQ) * sizeof(block_turbo3_tcq);
+    for (int64_t row = 0; row < nrows; row++) {
+        quantize_row_turbo3_tcq_ref(
+            src + row * n_per_row,
+            (block_turbo3_tcq *)((char *)dst + row * row_size),
+            n_per_row
+        );
+    }
+    return nrows * row_size;
+}
+
+void quantize_row_turbo2_tcq_ref(const float * GGML_RESTRICT x, block_turbo2_tcq * GGML_RESTRICT y, int64_t k) {
+    // Stub — CUDA kernel handles TCQ quantize (Viterbi). CPU path zeros out.
+    assert(k % QK_TURBO2_TCQ == 0);
+    const int nb = k / QK_TURBO2_TCQ;
+    for (int i = 0; i < nb; i++) {
+        float norm = 0.0f;
+        for (int j = 0; j < QK_TURBO2_TCQ; j++) norm += x[i*QK_TURBO2_TCQ + j] * x[i*QK_TURBO2_TCQ + j];
+        y[i].norm = GGML_FP32_TO_FP16(sqrtf(norm));
+        memset(y[i].qs, 0, 33);
+    }
+}
+
+void dequantize_row_turbo2_tcq(const block_turbo2_tcq * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) {
+    GGML_UNUSED(x);
+    assert(k % QK_TURBO2_TCQ == 0);
+    const int nb = k / QK_TURBO2_TCQ;
+    for (int block = 0; block < nb; block++) {
+        for (int j = 0; j < QK_TURBO2_TCQ; j++) {
+            y[block * QK_TURBO2_TCQ + j] = 0.0f;
+        }
+    }
+}
+
+size_t quantize_turbo2_tcq(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst,
+                         int64_t nrows, int64_t n_per_row, const float * imatrix) {
+    GGML_UNUSED(imatrix);
+    assert(n_per_row % QK_TURBO2_TCQ == 0);
+
+    size_t row_size = (n_per_row / QK_TURBO2_TCQ) * sizeof(block_turbo2_tcq);
+    for (int64_t row = 0; row < nrows; row++) {
+        quantize_row_turbo2_tcq_ref(
+            src + row * n_per_row,
+            (block_turbo2_tcq *)((char *)dst + row * row_size),
+            n_per_row
+        );
+    }
+    return nrows * row_size;
+}
+
 /* ---------- TURBO4_0: 3-bit PolarQuant + 1-bit QJL ---------- */
 
 void quantize_row_turbo4_0_ref(const float * GGML_RESTRICT x, block_turbo4_0 * GGML_RESTRICT y, int64_t k) {
diff --git a/llama.cpp/ggml/src/ggml.c b/llama.cpp/ggml/src/ggml.c
index d19e81f9b..2defaaefd 100644
--- a/llama.cpp/ggml/src/ggml.c
+++ b/llama.cpp/ggml/src/ggml.c
@@ -707,6 +707,32 @@ static const struct ggml_type_traits type_traits[GGML_TYPE_COUNT] = {
         .to_float                 = (ggml_to_float_t) dequantize_row_turbo8_0,
         .from_float_ref           = (ggml_from_float_t) quantize_row_turbo8_0_ref,
     },
+    // opencoti-hook: TCQ (F5 M6 C5) — Trellis-Coded Quantization KV (buun-aligned names)
+    [GGML_TYPE_TURBO3_TCQ] = {
+        .type_name                = "turbo3_tcq",
+        .blck_size                = QK_TURBO3_TCQ,
+        .type_size                = sizeof(block_turbo3_tcq),
+        .is_quantized             = true,
+        .to_float                 = (ggml_to_float_t) dequantize_row_turbo3_tcq,
+        .from_float_ref           = (ggml_from_float_t) quantize_row_turbo3_tcq_ref,
+    },
+    [GGML_TYPE_TURBO2_TCQ] = {
+        .type_name                = "turbo2_tcq",
+        .blck_size                = QK_TURBO2_TCQ,
+        .type_size                = sizeof(block_turbo2_tcq),
+        .is_quantized             = true,
+        .to_float                 = (ggml_to_float_t) dequantize_row_turbo2_tcq,
+        .from_float_ref           = (ggml_from_float_t) quantize_row_turbo2_tcq_ref,
+    },
+    // opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port) — scalar 6-bit KV tier, mirrors q5_0.
+    [GGML_TYPE_Q6_0] = {
+        .type_name                = "q6_0",
+        .blck_size                = QK6_0,
+        .type_size                = sizeof(block_q6_0),
+        .is_quantized             = true,
+        .to_float                 = (ggml_to_float_t) dequantize_row_q6_0,
+        .from_float_ref           = (ggml_from_float_t) quantize_row_q6_0_ref,
+    },
     [GGML_TYPE_TQ3_1S] = {
         .type_name                = "tq3_1s",
         .blck_size                = QK_TQ3_0,
@@ -8122,6 +8148,8 @@ size_t ggml_quantize_chunk(
         case GGML_TYPE_Q4_1:    result = quantize_q4_1   (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
         case GGML_TYPE_Q5_0:    result = quantize_q5_0   (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
         case GGML_TYPE_Q5_1:    result = quantize_q5_1   (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
+        // opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port) — see docs/features/poly_kv.md.
+        case GGML_TYPE_Q6_0:    result = quantize_q6_0   (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
         case GGML_TYPE_Q8_0:    result = quantize_q8_0   (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
         case GGML_TYPE_MXFP4:   result = quantize_mxfp4  (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
         case GGML_TYPE_NVFP4:   result = quantize_nvfp4  (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
@@ -8146,6 +8174,9 @@ size_t ggml_quantize_chunk(
         case GGML_TYPE_TURBO4_0: result = quantize_turbo4_0(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
         case GGML_TYPE_TURBO2_0: result = quantize_turbo2_0(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
         case GGML_TYPE_TURBO8_0: result = quantize_turbo8_0(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
+        // opencoti-hook: TCQ (F5 M6 C5)
+        case GGML_TYPE_TURBO3_TCQ: result = quantize_turbo3_tcq(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
+        case GGML_TYPE_TURBO2_TCQ: result = quantize_turbo2_tcq(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
         case GGML_TYPE_TQ3_1S:  result = quantize_tq3_1s(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
         case GGML_TYPE_TQ4_1S:  result = quantize_tq4_1s(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
         case GGML_TYPE_F16:
diff --git a/llama.cpp/ggml/src/iqk/iqk_common_extra.h b/llama.cpp/ggml/src/iqk/iqk_common_extra.h
index 969379271..85f006006 100644
--- a/llama.cpp/ggml/src/iqk/iqk_common_extra.h
+++ b/llama.cpp/ggml/src/iqk/iqk_common_extra.h
@@ -14,13 +14,11 @@ typedef struct {
 } block_q5_0_r4;
 static_assert(sizeof(block_q5_0_r4) == 4*sizeof(ggml_half) + QK5_0*2 + QK5_0/2, "wrong q5_0_r4 block size/padding");
 
-#define QK6_0 32
-typedef struct {
-    ggml_half d;         // delta
-    uint8_t qh[QK6_0/4]; // 5+6-th bit of quants
-    uint8_t qs[QK6_0/2]; // nibbles / quants
-} block_q6_0;
-static_assert(sizeof(block_q6_0) == sizeof(ggml_half) + QK6_0/2 + QK6_0/4, "wrong q6_0 block size/padding");
+#define QK6_0 32   // kept: block_q6_0_r4 below uses it; identical to ggml-common.h's #define (harmless redef)
+// opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port, F-A) — block_q6_0 now lives in the real ggml-common.h
+// (this delta's shim includes "../ggml-common.h" BEFORE this file), so the ik copy is removed. Two anonymous
+// -struct typedefs of the same name are DISTINCT types in C++ → "typedef redefinition with different types"
+// (bug-611). Re-sync per this header's contract: drop block types once the main ggml-common.h provides them.
 
 typedef struct {
     ggml_half d[4];      // delta
diff --git a/llama.cpp/ggml/src/iqk/iqk_ggml_type_ext.h b/llama.cpp/ggml/src/iqk/iqk_ggml_type_ext.h
index 673828191..0b1d97353 100644
--- a/llama.cpp/ggml/src/iqk/iqk_ggml_type_ext.h
+++ b/llama.cpp/ggml/src/iqk/iqk_ggml_type_ext.h
@@ -30,9 +30,10 @@
 #ifndef GGML_TYPE_Q8_2_X4
 #define GGML_TYPE_Q8_2_X4 ((ggml_type)99)
 #endif
-#ifndef GGML_TYPE_Q6_0
-#define GGML_TYPE_Q6_0 ((ggml_type)133)
-#endif
+// opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port, F-A) — GGML_TYPE_Q6_0 is now a NATIVE llamafile enum
+// value (=50). The old ik delta `#define GGML_TYPE_Q6_0 ((ggml_type)133)` is removed: #ifndef cannot see an
+// enum constant (only macros), so it would define 133 and SHADOW the real 50 → iqk and main ggml disagree on
+// q6_0's identity (bug-611). iqk TUs already see ggml.h's enum (they use Q4_0/Q8_0/etc.), so 50 resolves cleanly.
 #ifndef GGML_TYPE_IQ1_BN
 #define GGML_TYPE_IQ1_BN ((ggml_type)134)
 #endif
diff --git a/llama.cpp/src/dca.cpp b/llama.cpp/src/dca.cpp
index 78e0b43bc..b84ff8374 100644
--- a/llama.cpp/src/dca.cpp
+++ b/llama.cpp/src/dca.cpp
@@ -175,6 +175,35 @@ ggml_tensor * llm_graph_context::build_dca_rope(ggml_tensor * x, ggml_tensor * p
             ext_factor, attn_factor, beta_fast, beta_slow);
 }
 
+// opencoti F5 dca #444 (DCA all-KV C1): dequant-on-lift a cached K/V tensor to f16 for the f16-only
+// fused DCA kernel (the cache itself stays quantized — that is the memory win; we only materialize a
+// transient f16 copy for the kernel). The CUDA CPY backend has DIRECT ->f16 casts for f32/bf16/turbo
+// (turbo via the #385 InnerQ-aware kernel), and scalar-quant->f32 casts (q8_0/q4_0/q4_1/q5_0/q5_1),
+// but NO direct scalar-quant->f16 — that pair fails CPY supports_op and the scheduler CPU-spills it
+// (whole-cache dequant on the CPU + H2D copy per layer per token = a decode cliff). So route scalar-
+// quantized caches through f32 first: both legs are GPU CPY kernels, and the f32 intermediate is
+// NUMERICALLY IDENTICAL to a one-pass dequant (d*q is computed in f32 then rounded to f16 either way),
+// so this is byte-exact to the eventual single-pass quant->f16 kernel. f16 returns unchanged (byte-
+// identical to the prior assert path). The single-pass direct quant->f16 CPY kernels are the C-series
+// perf follow-on (built + restamped at #448); this helper keeps C1 host-only and immediately gateable.
+static ggml_tensor * dca_lift_to_f16(ggml_context * ctx0, ggml_tensor * x) {
+    if (x->type == GGML_TYPE_F16) {
+        return x;
+    }
+    const bool scalar_quant =
+        x->type == GGML_TYPE_Q8_0 || x->type == GGML_TYPE_Q4_0 || x->type == GGML_TYPE_Q4_1 ||
+        x->type == GGML_TYPE_Q5_0 || x->type == GGML_TYPE_Q5_1 ||
+        // opencoti-hook: DCA all-KV #444 — q6_0 (added later, #512) lifts via f32 like the other
+        // scalar quants. Omitting it here sent q6_0 through the direct ggml_cast(q6_0->f16) below,
+        // which has no CUDA cpy kernel (only q6_0<->f32 is wired) -> CPU dup_from_q abort at
+        // ggml-cpu/ops.cpp:594. q6_0<->f32 cpy IS GPU-supported, so the two-leg lift stays on GPU.
+        x->type == GGML_TYPE_Q6_0;
+    if (scalar_quant) {
+        x = ggml_cast(ctx0, x, GGML_TYPE_F32);
+    }
+    return ggml_cast(ctx0, x, GGML_TYPE_F16);
+}
+
 ggml_tensor * llm_graph_context::build_attn_dca_core(
         const llama_kv_cache_context * mctx_cur,
         llm_graph_input_dca * inp_dca,
@@ -190,12 +219,24 @@ ggml_tensor * llm_graph_context::build_attn_dca_core(
     //     kept correct under context-shift by the DCA-aware modular K-shift. So read it DIRECTLY — no
     //     per-step full-cache rope. That re-rope of the entire cache every graph eval was the
     //     O(n_kv)/token decode bottleneck (gen halving per context-doubling). get_k returns
-    //     [n_embd_head, n_head_kv, n_kv, 1]. DCA requires an f16 K cache (no q8_0/turbo on this path).
-    ggml_tensor * k_dca = mctx_cur->get_k(ctx0, il);
-    GGML_ASSERT(k_dca->type == GGML_TYPE_F16 && "dca: requires an f16 K cache (no q8_0/turbo KV)");
+    //     [n_embd_head, n_head_kv, n_kv, 1].
+    //     opencoti F5 dca #444 (DCA all-KV C1): the fused DCA kernel is f16-only by construction, so a
+    //     quantized/bf16/turbo K cache is DEQUANTIZED-ON-LIFT to a transient f16 here (the cache itself
+    //     stays quantized — that is the memory win; dca_lift_to_f16 only materializes an f16 copy for
+    //     the kernel). At f16 this is a no-op (byte-identical to the prior assert path). Context-shift
+    //     stays correct independently: the modular K-shift re-ropes the STORED, still-quantized cache
+    //     via build_rope_shift's existing dequant->rope->requant quantized branch (llama-kv-cache.cpp).
+    ggml_tensor * k_dca = dca_lift_to_f16(ctx0, mctx_cur->get_k(ctx0, il));
 
     // (2) read V once; permute K + V to the flash-attn layout [n_embd_head, seq, n_head, n_stream].
+    //     opencoti F5 dca #444 (DCA all-KV C1): lift a block/scalar-quantized V (q8_0/turbo/...) to f16
+    //     while it is still CONTIGUOUS — a permuted view of a block type is not cleanly cpy-able (the
+    //     block stride spans dim-0), mirroring build_attn_mha's pre-permute turbo cast. f32 V keeps the
+    //     post-permute cast below (unchanged); f16 V is untouched (byte-identical).
     ggml_tensor * v = mctx_cur->get_v(ctx0, il);
+    if (v->type != GGML_TYPE_F16 && v->type != GGML_TYPE_F32) {
+        v = dca_lift_to_f16(ctx0, v);
+    }
     const bool v_trans = v->nb[1] > v->nb[2];
     ggml_tensor * vp = ggml_permute(ctx0, v, 0, 2, 1, 3);
     if (v_trans) {
diff --git a/llama.cpp/src/llama-graph.cpp b/llama.cpp/src/llama-graph.cpp
index 92e3f0b48..2fe40de6f 100644
--- a/llama.cpp/src/llama-graph.cpp
+++ b/llama.cpp/src/llama-graph.cpp
@@ -2066,7 +2066,11 @@ ggml_tensor * llm_graph_context::build_attn_mha(
     // those still cast turbo→f16 up front (ggml_cuda_cpy → ggml_cpy_turbo_f16_cuda, WHT-once +
     // InnerQ, VRAM-resident). A single Gemma model thus runs fused SWA layers + materialized global.
     const auto is_turbo_vec = [](enum ggml_type t) {
-        return t == GGML_TYPE_TURBO2_0 || t == GGML_TYPE_TURBO3_0 || t == GGML_TYPE_TURBO4_0;
+        return t == GGML_TYPE_TURBO2_0 || t == GGML_TYPE_TURBO3_0 || t == GGML_TYPE_TURBO4_0
+            // opencoti-hook: TCQ FA-VEC decode (#447/#500). Trellis turbo joins the fused FA-VEC
+            // predicate: K/V are FWHT-rotated (no InnerQ), so the same graph Q-rotation (ggml_turbo_wht
+            // dir=0, scale=nullptr → pure FWHT since InnerQ is inactive) + output inverse-WHT apply.
+            || t == GGML_TYPE_TURBO3_TCQ || t == GGML_TYPE_TURBO2_TCQ;
     };
     // opencoti-hook: turboquant perf-lift (M6-S2 Level B, L4.5 AB-prefill-hybrid). The fused FA-VEC
     // turbo read wins at DECODE (in-register, bandwidth-bound: turbo2/3/4 = 56/51/48 tps vs 43
@@ -2088,12 +2092,26 @@ ggml_tensor * llm_graph_context::build_attn_mha(
     // scales with n_kv). turbo4/turbo8 at 512 still take the materialize fallback (no D=512 instance).
     // The WHT is block-diagonal over 128-blocks, so 512 = 4 Parseval-exact 128-blocks → exact.
     const bool fused_turbo_512 =
-        k->ne[0] == 512 && (k->type == GGML_TYPE_TURBO2_0 || k->type == GGML_TYPE_TURBO3_0);
+        k->ne[0] == 512 && (k->type == GGML_TYPE_TURBO2_0 || k->type == GGML_TYPE_TURBO3_0
+            || k->type == GGML_TYPE_TURBO3_TCQ || k->type == GGML_TYPE_TURBO2_TCQ); // opencoti TCQ #447/#500 (Gemma-4 GLOBAL)
+    // opencoti-hook: TCQ FA-VEC decode (#447/#500). TCQ has NO materialize-lift (turbo-cpy.cu has no
+    // TCQ→f16 inverse-FWHT cast yet), so it must ALWAYS fuse — prefill (large n_q) rides the VEC-tiled
+    // path (correct, slower than MMA; a prefill-MMA dequant is a deferred perf follow-on, like buun's
+    // k_turbo3_tcq_dequant_f16). Without this, a prefill ubatch would fall to the else-branch with no
+    // cast and emit raw TCQ to FA with an UN-rotated Q → wrong logits.
+    const bool is_tcq_kv = (k->type == GGML_TYPE_TURBO3_TCQ || k->type == GGML_TYPE_TURBO2_TCQ); // opencoti-hook: TCQ FA-VEC (#447/#500) — no materialize-lift, always fuse
+    // opencoti-hook: asymmetric TCQ (#513). TCQ has no f16 materialize-lift, so the only correct path is
+    // the in-register fused FA-VEC read. Allow K!=V when BOTH sides are TCQ (e.g. turbo3_tcq K / turbo2_tcq
+    // V): K and V are FWHT-rotated identically, so the same Q-rotation + output inverse-WHT stay logit-exact
+    // regardless of per-side trellis bit-width. Backed by the dedicated asymmetric FA-VEC instances
+    // (fattn-vec-instance-turbo3_tcq-turbo2_tcq + reverse); the CUDA selector K!=V gate is relaxed to match.
+    const bool both_tcq = (k->type == GGML_TYPE_TURBO3_TCQ || k->type == GGML_TYPE_TURBO2_TCQ) &&
+                          (v->type == GGML_TYPE_TURBO3_TCQ || v->type == GGML_TYPE_TURBO2_TCQ);
     const bool fused_turbo =
         cparams.flash_attn && kq_b == nullptr &&
-        is_turbo_vec(k->type) && k->type == v->type &&
+        is_turbo_vec(k->type) && is_turbo_vec(v->type) && (k->type == v->type || both_tcq) &&
         (k->ne[0] == 128 || k->ne[0] == 256 || fused_turbo_512) &&
-        tbv_nq_per_strm <= tbv_fuse_max_nq;
+        (tbv_nq_per_strm <= tbv_fuse_max_nq || is_tcq_kv);
 
     if (fused_turbo) {
         // forward-rotate Q (×scale_inv → signs₁ → WHT → 1/√G·signs₂); group auto-picks 128 since
@@ -2856,7 +2874,8 @@ ggml_tensor * llm_graph_context::build_attn_mtp(
     ggml_tensor * k = mctx_cur->get_k(ctx0, il_kv_tgt);
     ggml_tensor * v = use_k_as_v ? k : mctx_cur->get_v(ctx0, il_kv_tgt);
 
-    if (k->type == GGML_TYPE_TURBO3_0 || k->type == GGML_TYPE_TURBO4_0 || k->type == GGML_TYPE_TURBO2_0) {
+    if (k->type == GGML_TYPE_TURBO3_0 || k->type == GGML_TYPE_TURBO4_0 || k->type == GGML_TYPE_TURBO2_0
+        || k->type == GGML_TYPE_TURBO3_TCQ || k->type == GGML_TYPE_TURBO2_TCQ) { // opencoti-hook: TCQ (#447/#500)
         // opencoti F5 M6-S4 mtp (P3 R1 fix): do NOT pre-rotate Q here. build_attn_mha owns the
         // authoritative fused-turbo path — it forward-rotates Q (ggml_turbo_wht dir=0, graph.cpp
         // :1950) AND applies the paired output inverse-WHT (dir=1, :2001) for turbo2/3/4 K==V at
@@ -2875,7 +2894,8 @@ ggml_tensor * llm_graph_context::build_attn_mtp(
     ggml_tensor * cur = build_attn_mha(q, k, v, kq_b, kq_mask, sinks, v_mla, kq_scale, il_mtp);
     cb(cur, "kqv_out_mtp", il_mtp);
 
-    if (v->type == GGML_TYPE_TURBO3_0 || v->type == GGML_TYPE_TURBO4_0 || v->type == GGML_TYPE_TURBO2_0) {
+    if (v->type == GGML_TYPE_TURBO3_0 || v->type == GGML_TYPE_TURBO4_0 || v->type == GGML_TYPE_TURBO2_0
+        || v->type == GGML_TYPE_TURBO3_TCQ || v->type == GGML_TYPE_TURBO2_TCQ) { // opencoti-hook: TCQ (#447/#500)
         const int64_t orig_v_head   = kv_embd_head_v;
         const int64_t padded_v_head = v->ne[0];
         if (padded_v_head != orig_v_head) {
diff --git a/llama.cpp/tools/perplexity/main.cpp b/llama.cpp/tools/perplexity/main.cpp
index 13a9940e9..0f1e79fef 100644
--- a/llama.cpp/tools/perplexity/main.cpp
+++ b/llama.cpp/tools/perplexity/main.cpp
@@ -1,5 +1,14 @@
+// opencoti (#517 F-E): trigger llamafile's GPU bootstrap (dlopen CUDA DSO + ggml_backend_register)
+// before the upstream perplexity entry, mirroring tools/server/server.cpp main()'s llamafile_has_gpu()
+// call. The 0.10.3 bump replaced llamafile's patched perplexity main with this thin upstream wrapper,
+// which never registers the llamafile CUDA backend -> the tool runs CPU-only ("no usable GPU found")
+// and CUDA-only turbo/TCQ ops (TURBO_WHT) abort. FLAG_gpu defaults to AUTO, so the detect call alone
+// loads+registers the backend (same as the server). Eval-only tool; see docs §5a-TCQ / buglog.
+extern "C" bool llamafile_has_gpu(void);
+
 int llama_perplexity(int argc, char ** argv);
 
 int main(int argc, char ** argv) {
+    llamafile_has_gpu();
     return llama_perplexity(argc, argv);
 }
diff --git a/llamafile/build-functions.sh b/llamafile/build-functions.sh
index 87279b8..4781516 100644
--- a/llamafile/build-functions.sh
+++ b/llamafile/build-functions.sh
@@ -154,7 +154,39 @@ collect_gpu_sources() {
     if [ "$fa_all_quants" != "1" ]; then
         for f in "$ti_dir"/fattn-vec-instance-turbo2_0-turbo2_0.cu \
                  "$ti_dir"/fattn-vec-instance-turbo3_0-turbo3_0.cu \
-                 "$ti_dir"/fattn-vec-instance-turbo4_0-turbo4_0.cu; do
+                 "$ti_dir"/fattn-vec-instance-turbo4_0-turbo4_0.cu \
+                 "$ti_dir"/fattn-vec-instance-turbo3_tcq-turbo3_tcq.cu \
+                 "$ti_dir"/fattn-vec-instance-turbo2_tcq-turbo2_tcq.cu; do
+            [ -f "$f" ] && CUDA_SOURCES="$CUDA_SOURCES $f"
+        done
+    fi
+
+    # 3c. opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port). q6_0 is a scalar KV type read by the
+    #     FA-VEC path; its 5 Anbeeld instances must compile in the DEFAULT build (same #468/bug-339
+    #     reasoning as the turbo block above — the default build does NOT glob fattn-vec, so without
+    #     this the DSO lacks ggml_cuda_flash_attn_ext_vec_case<D,q6_0,*> and dlopen fails the instant
+    #     q6_0 KV reaches flash-attention). No-op under FA_ALL_QUANTS=1 (already globbed) and on a
+    #     vanilla tree (per-file [ -f ] guard).
+    if [ "$fa_all_quants" != "1" ]; then
+        for f in "$ti_dir"/fattn-vec-instance-q6_0-q6_0.cu \
+                 "$ti_dir"/fattn-vec-instance-q8_0-q6_0.cu \
+                 "$ti_dir"/fattn-vec-instance-q6_0-q5_0.cu \
+                 "$ti_dir"/fattn-vec-instance-q6_0-f16.cu \
+                 "$ti_dir"/fattn-vec-instance-f16-q6_0.cu; do
+            [ -f "$f" ] && CUDA_SOURCES="$CUDA_SOURCES $f"
+        done
+    fi
+
+    # 3d. opencoti-hook: asymmetric TCQ (#513). The mixed turbo3_tcq/turbo2_tcq FA-VEC instances (both
+    #     directions) read TCQ K from one tier and TCQ V from the other; like the symmetric TCQ block (3b)
+    #     these are FUSED in-register (no dequant-on-lift inverse for TCQ exists), so they MUST compile in
+    #     the DEFAULT build — without them the DSO lacks ggml_cuda_flash_attn_ext_vec_case<D,turbo3_tcq,
+    #     turbo2_tcq> (and reverse) and dlopen fails the instant asymmetric TCQ KV reaches flash-attention
+    #     (same #468/bug-339 reasoning). No-op under FA_ALL_QUANTS=1 (already globbed) and on a vanilla
+    #     tree (per-file [ -f ] guard).
+    if [ "$fa_all_quants" != "1" ]; then
+        for f in "$ti_dir"/fattn-vec-instance-turbo3_tcq-turbo2_tcq.cu \
+                 "$ti_dir"/fattn-vec-instance-turbo2_tcq-turbo3_tcq.cu; do
             [ -f "$f" ] && CUDA_SOURCES="$CUDA_SOURCES $f"
         done
     fi