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 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 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 #include #include 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; + case GGML_TYPE_Q6_0: // opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port) + return dequantize_block_cont_cuda; case GGML_TYPE_Q5_1: return dequantize_block_cont_cuda; 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; + case GGML_TYPE_Q6_0: // opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port) + return dequantize_block_cont_cuda; case GGML_TYPE_Q5_1: return dequantize_block_cont_cuda; 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; case GGML_TYPE_Q5_0: return dequantize_block_cuda; + case GGML_TYPE_Q6_0: // opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port) + return dequantize_block_cuda; case GGML_TYPE_Q5_1: return dequantize_block_cuda; 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; case GGML_TYPE_Q5_0: return dequantize_block_cuda; + case GGML_TYPE_Q6_0: // opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port) + return dequantize_block_cuda; case GGML_TYPE_Q5_1: return dequantize_block_cuda; 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; case GGML_TYPE_Q5_0: return dequantize_block_cuda; + case GGML_TYPE_Q6_0: // opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port) + return dequantize_block_cuda; case GGML_TYPE_Q5_1: return dequantize_block_cuda; 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<<>> + (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, QK6_0><<>>( + 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. +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 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 +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(&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(&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 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 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 +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(&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) { + 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) { + 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, "bad type"); + } +} + template 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 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; } else if constexpr (type_K == GGML_TYPE_TURBO4_0) { return vec_dot_fattn_vec_KQ_turbo4; + } 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; + } else if constexpr (type_K == GGML_TYPE_TURBO2_TCQ) { + return vec_dot_fattn_vec_KQ_turbo2_tcq; } else if constexpr (type_K == GGML_TYPE_Q4_0) { return vec_dot_fattn_vec_KQ_q4_0; } else if constexpr (type_K == GGML_TYPE_Q4_1) { return vec_dot_fattn_vec_KQ_q4_1; } else if constexpr (type_K == GGML_TYPE_Q5_0) { return vec_dot_fattn_vec_KQ_q5_0; + } 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; } else if constexpr (type_K == GGML_TYPE_Q5_1) { return vec_dot_fattn_vec_KQ_q5_1; } 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; } else if constexpr (type_V == GGML_TYPE_TURBO4_0) { return dequantize_V_turbo4; + } else if constexpr (type_V == GGML_TYPE_TURBO3_TCQ) { // opencoti-hook: TCQ FA-VEC decode (#447/#500) + return dequantize_V_turbo3_tcq; + } else if constexpr (type_V == GGML_TYPE_TURBO2_TCQ) { + return dequantize_V_turbo2_tcq; } else if constexpr (type_V == GGML_TYPE_Q4_0) { return dequantize_V_q4_0; } else if constexpr (type_V == GGML_TYPE_Q4_1) { return dequantize_V_q4_1; } else if constexpr (type_V == GGML_TYPE_Q5_0) { return dequantize_V_q5_0; + } else if constexpr (type_V == GGML_TYPE_Q6_0) { // opencoti-hook: q6_0 KV (Anbeeld/beellama.cpp port) + return dequantize_V_q6_0; } else if constexpr (type_V == GGML_TYPE_Q5_1) { return dequantize_V_q5_1; } 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 +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 +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(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 +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 +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(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 +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) { + 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) { + for (int l = 0; l < ne; ++l) ((float *)dst)[l] = vals[l]; + } else { static_assert(std::is_same_v, "bad type"); } +} + +// Wrapper using __constant__ codebook (for function pointer dispatch via dequantize_V_t) +template +static __device__ __forceinline__ void dequantize_V_turbo3_tcq( + const void * __restrict__ vx, void * __restrict__ dst, const int64_t i0) { + dequantize_V_turbo3_tcq_cb(vx, dst, i0, d_turbo3_tcq_codebook_fattn); +} + +// TCQ 2-bit V dequant: 8-bit state → codebook lookup +template +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) { + 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) { + for (int l = 0; l < ne; ++l) ((float *)dst)[l] = vals[l]; + } else { static_assert(std::is_same_v, "bad type"); } +} + +// Wrapper using __constant__ codebook (for function pointer dispatch via dequantize_V_t) +template +static __device__ __forceinline__ void dequantize_V_turbo2_tcq( + const void * __restrict__ vx, void * __restrict__ dst, const int64_t i0) { + dequantize_V_turbo2_tcq_cb(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 // strncmp (cache_k_ vs cache_v_ detection for TCQ alpha) #include #include @@ -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 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= +// 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 — 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 +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), tcq3_bt_shared_bytes); + else + CUDA_SET_SHARED_MEMORY_LIMIT(k_set_rows_turbo3_tcq, 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<<<(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<<<(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 +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), tcq2_bt_shared_bytes); + else + CUDA_SET_SHARED_MEMORY_LIMIT(k_set_rows_turbo2_tcq, 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<<<(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<<<(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 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( + 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( 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(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(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(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 +// launch resolves to the float instantiation, byte-identical to buun. +// The host dispatches the 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 +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 +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 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 (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