From: opencoti Subject: [PATCH 0073] F5 M6 PolyKV S2 — TurboQuant KV tier family + fused FA-VEC TurboQuant KV (F5 M6 S2, #371/#374-396): a non-uniform Lloyd-Max centroid + Walsh-Hadamard-rotation KV-cache quant ladder TURBO{2,3,4,8}_0 (block=group=128) with per-channel InnerQ K-variance equalization. Three layers ship together: - tiers + InnerQ (#371/#374-382): ggml type registration, block structs + centroids, CPU quant/dequant (ggml-turbo-quant.c), GPU set-rows writers, the device-symbol InnerQ coordinator (calibrate/apply, B-adapted per head_dim). - Level-A materialize (#383-388): an InnerQ-aware turbo->f16 GPU lift on the M7 dequant-on-lift path so any tier runs through stock f16 FA. - Level-B fused FA-VEC (#390-396): a new op GGML_OP_TURBO_WHT rotates the small n_q-sized Q (fwd) and FA output (inv) so flash-attention reads turbo2/3/4 K/V IN-REGISTER (centroid x norm, WHT-rotated domain, no f16 materialize). Parseval (WHT(Q/s).WHT(K*s)=Q.K) cancels InnerQ on the K/V read. build_attn_mha fuses decode (n_q<=TBV_FUSE_MAX_NQ, default 8) at head_dim in {128,256}; turbo8 + head_dim-512 keep the Level-A materialize fallback. Acceptance = LOGIT-EQUIVALENCE vs f16 (#355 harness), NEVER greedy byte-equality (turboN is lossy). 4/4 tiers PASS; turbo3 mean_tv 0.0067 (tighter than q4_0's 0.0094) at 90 vs 129 MiB; fused decode scales with compression (turbo2 56 > turbo3 51 > turbo4 48 tps @6k, breaking the Level-A flat ~41 ceiling). Composes on the full F5 stack (gate .opencoti/m6-s3-glue-gate.sh). Bench: .opencoti/s2-tier-family-gate.sh, s2-tier-overview.sh, m6-s3-glue-gate.sh Milestone: F5 M6 PolyKV S2 — see docs/features/poly_kv.md ABI: GGML_OP_COUNT bumped by GGML_OP_TURBO_WHT (appended at end); cosmocc binary + ggml-cuda.so DSO rebuilt and paired. Applies after: 0071 (numbered 0073; the TurboQuant edits to fattn.cu / llama-graph.cpp / ggml-cuda.cu sit ON TOP of 0070's rewrite of those files, so lexical apply order MUST land this after 0070/0071; the original 0061 slot predates 0070 being built first). diff --git a/llama.cpp/BUILD.mk b/llama.cpp/BUILD.mk --- a/llama.cpp/BUILD.mk +++ b/llama.cpp/BUILD.mk @@ -17,6 +17,7 @@ LLAMA_COMMIT := $(shell cd llama.cpp 2>/dev/null && git rev-parse --short HEAD 2 GGML_SRCS_C := \ llama.cpp/ggml/src/ggml-alloc.c \ + llama.cpp/ggml/src/ggml-turbo-quant.c \ llama.cpp/ggml/src/ggml-quants.c \ llama.cpp/ggml/src/ggml.c \ llama.cpp/ggml/src/ggml-cpu/ggml-cpu.c \ diff --git a/llama.cpp/common/arg.cpp b/llama.cpp/common/arg.cpp --- a/llama.cpp/common/arg.cpp +++ b/llama.cpp/common/arg.cpp @@ -412,6 +412,11 @@ const std::vector kv_cache_types = { GGML_TYPE_IQ4_NL, GGML_TYPE_Q5_0, GGML_TYPE_Q5_1, + // 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, }; 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 --- a/llama.cpp/ggml/include/ggml.h +++ b/llama.cpp/ggml/include/ggml.h @@ -429,7 +429,15 @@ extern "C" { GGML_TYPE_MXFP4 = 39, // MXFP4 (1 block) GGML_TYPE_NVFP4 = 40, // NVFP4 (4 blocks, E4M3 scale) GGML_TYPE_Q1_0 = 41, - GGML_TYPE_COUNT = 42, + // opencoti-hook: turboquant — see docs/features/poly_kv.md (M6-S2, AtomicBot fork adopt) + GGML_TYPE_TURBO2_0 = 42, // TurboQuant 2-bit KV: WHT + 2-bit PolarQuant + GGML_TYPE_TURBO3_0 = 43, // TurboQuant 3-bit KV: WHT + 3-bit PolarQuant + GGML_TYPE_TURBO4_0 = 44, // TurboQuant 4-bit KV: WHT + 4-bit PolarQuant + GGML_TYPE_TQ3_1S = 45, // TurboQuant 3-bit weight: WHT-rotated Lloyd-Max, block=32 + 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, }; // precision @@ -589,6 +597,8 @@ extern "C" { GGML_OP_FLASH_ATTN_TAIL_PARTIAL, // opencoti-hook: f5-rolling-kv — W4 M7 S3-a #353 + GGML_OP_TURBO_WHT, // opencoti-hook: turboquant perf-lift (Level B) — F5 M6-S2 #391 + GGML_OP_COUNT, }; @@ -2481,6 +2491,21 @@ extern "C" { struct ggml_tensor * a, enum ggml_prec prec); + // opencoti-hook: turboquant perf-lift (Level B) — F5 M6-S2 #391. Walsh–Hadamard + // rotation op for the fused turbo FA-VEC path. direction 0 = forward (Q + // pre-rotation: [InnerQ scale_inv] → signs1 → WHT → 1/√G·signs2); direction 1 = + // inverse (FA-output un-rotation: signs2 → WHT → 1/√G·signs1 → [InnerQ scale_inv]). + // group_size 0 ⇒ auto (128 if ne[0]%128==0 else 64). InnerQ scale is B-adapted: a + // is f32; pass scale=NULL and the CUDA kernel resolves InnerQ from the device + // symbol (turbo_innerq_active_device_scale, keyed on ne[0]); a non-NULL scale src + // (fork-compat) takes precedence. See docs/features/poly_kv.md. + GGML_API struct ggml_tensor * ggml_turbo_wht( + struct ggml_context * ctx, + struct ggml_tensor * a, + int direction, + int group_size, + struct ggml_tensor * scale); + GGML_API enum ggml_prec ggml_flash_attn_ext_get_prec( const struct ggml_tensor * a); diff --git a/llama.cpp/ggml/src/ggml-backend.cpp b/llama.cpp/ggml/src/ggml-backend.cpp --- a/llama.cpp/ggml/src/ggml-backend.cpp +++ b/llama.cpp/ggml/src/ggml-backend.cpp @@ -903,6 +903,19 @@ static int ggml_backend_sched_backend_id_from_cur(ggml_backend_sched_t sched, st if (tensor->buffer || (tensor->view_src && tensor->view_src->buffer)) { // since the tensor is pre-allocated, it cannot be moved to another backend ggml_backend_buffer_t buffer = tensor->view_src ? tensor->view_src->buffer : tensor->buffer; + // OPENCOTI-DEBUG-TURBO3 (temporary, M6-S2 Gate-2): dump why no backend matched + fprintf(stderr, "[OPENCOTI-DBG] node=%s op=%s op->type=%s(%d) src0=%s src1=%s src2=%s\n", + tensor->name, ggml_op_name(tensor->op), + ggml_type_name(tensor->type), (int)tensor->type, + tensor->src[0] ? ggml_type_name(tensor->src[0]->type) : "null", + tensor->src[1] ? ggml_type_name(tensor->src[1]->type) : "null", + tensor->src[2] ? ggml_type_name(tensor->src[2]->type) : "null"); + for (int bi = 0; bi < sched->n_backends; bi++) { + const bool sb = ggml_backend_supports_buft(sched->backends[bi], buffer->buft); + const bool so = ggml_backend_supports_op(sched->backends[bi], tensor); + fprintf(stderr, "[OPENCOTI-DBG] backend[%d]=%s buft=%s supports_buft=%d supports_op=%d\n", + bi, ggml_backend_name(sched->backends[bi]), ggml_backend_buft_name(buffer->buft), sb, so); + } GGML_ABORT("pre-allocated tensor (%s) in a buffer (%s) that cannot run the operation (%s)", tensor->name, ggml_backend_buffer_name(buffer), ggml_op_name(tensor->op)); } diff --git a/llama.cpp/ggml/src/ggml-common.h b/llama.cpp/ggml/src/ggml-common.h --- a/llama.cpp/ggml/src/ggml-common.h +++ b/llama.cpp/ggml/src/ggml-common.h @@ -449,6 +449,113 @@ typedef struct { } block_iq4_xs; static_assert(sizeof(block_iq4_xs) == sizeof(ggml_half) + sizeof(uint16_t) + QK_K/64 + QK_K/2, "wrong iq4_xs block size/padding"); +// opencoti-hook: turboquant — see docs/features/poly_kv.md (M6-S2) +// Per block: norm(fp16) + 2-bit indices (8 bytes) + 1-bit extra (4 bytes) = 14 bytes per 32 values +// = 3.5 bits/value → 4.6× compression vs fp16 +// The 3-bit index is split: lower 2 bits in qs[], upper 1 bit in signs[] +#define QK_TURBO3 128 // Block size 128: one block per rotation group, eliminates redundant norms +#define QK_TURBO3_GROUP 128 // rotation group size = head_dim +#define QR_TURBO3 1 // opencoti-hook: turboquant (M6-S2) — dequant call emits 2 consecutive elements (float2) +// Derived: FA template nl parameters (auto-scale with block size) +#define NL_TURBO3 (QK_TURBO3 / 16) // non-vec FA iterations per block +#define NL_TURBO3_VEC (QK_TURBO3 / 4) // vec FA iterations per block +typedef struct { + ggml_half norm; // 2 bytes: vector L2 norm (for rescaling) + uint8_t qs[QK_TURBO3 / 4]; // 8 bytes: lower 2-bit indices (4 per byte) + uint8_t signs[QK_TURBO3 / 8]; // 4 bytes: upper 1-bit of 3-bit index (8 per byte) +} block_turbo3_0; // 14 bytes total +static_assert(sizeof(block_turbo3_0) == sizeof(ggml_half) + QK_TURBO3/4 + QK_TURBO3/8, "wrong turbo3_0 block size/padding"); + +// TurboQuant 4-bit: 3-bit PolarQuant indices + 1-bit QJL signs +// TURBO4_USE_4BIT: switch between 4-bit PolarQuant (new) and 3-bit+QJL (legacy) +// Default: 4-bit on all backends (Metal + CUDA validated) +#ifndef TURBO4_USE_4BIT +# define TURBO4_USE_4BIT 1 +#endif + +#define QK_TURBO4 128 +#define QR_TURBO4 1 // opencoti-hook: turboquant (M6-S2 tier family) — float2 dequant, same as turbo3 + +#if TURBO4_USE_4BIT +// 4-bit PolarQuant: 16 optimal centroids, nibble packed, no QJL +// Per block: norm(fp16) + rnorm(fp16, reserved) + 4-bit indices (64 bytes) +// = 68 bytes per 128 values = 4.25 bits/value → 3.8× compression vs fp16 +typedef struct { + ggml_half norm; // 2 bytes + ggml_half rnorm; // 2 bytes (reserved, unused in 4-bit mode) + uint8_t qs[QK_TURBO4 / 2]; // 64 bytes: 4-bit PolarQuant indices (nibble packed) +} block_turbo4_0; // 68 bytes total +static_assert(sizeof(block_turbo4_0) == 68, "wrong turbo4_0 block size"); +#else +// Legacy 3-bit PolarQuant + 1-bit QJL (original paper design) +// Per block: norm(fp16) + rnorm(fp16) + 3-bit indices (48 bytes) + 1-bit QJL signs (16 bytes) +// = 68 bytes per 128 values = 4.25 bits/value → 3.8× compression vs fp16 +typedef struct { + ggml_half norm; // 2 bytes + ggml_half rnorm; // 2 bytes: residual norm for QJL scale + uint8_t qs[QK_TURBO4 * 3 / 8]; // 48 bytes: 3-bit PolarQuant indices + uint8_t signs[QK_TURBO4 / 8]; // 16 bytes: 1-bit QJL signs +} block_turbo4_0; // 68 bytes total +static_assert(sizeof(block_turbo4_0) == 2*sizeof(ggml_half) + QK_TURBO4*3/8 + QK_TURBO4/8, "wrong turbo4_0 block size"); +#endif + +static_assert(QK_TURBO4 == 128, "turbo4 kernels assume QK_TURBO4 == 128"); + +// TurboQuant 2-bit: 2-bit PolarQuant indices only (no QJL) +// Per block: norm(fp16) + 2-bit indices (8 bytes) = 10 bytes per 32 values +// = 2.5 bits/value → 6.4× compression vs fp16 +// 4 centroids (Lloyd-Max for N(0, 1/128)): {-0.133462, -0.039994, 0.039994, 0.133462} +#define QK_TURBO2 128 // Block size 128: one block per rotation group +#define QK_TURBO2_GROUP 128 // rotation group size = head_dim +#define QR_TURBO2 1 // opencoti-hook: turboquant (M6-S2 tier family) — float2 dequant, same as turbo3 +// Derived: FA template nl parameters (auto-scale with block size) +#define NL_TURBO2 (QK_TURBO2 / 16) // non-vec FA iterations per block +#define NL_TURBO2_VEC (QK_TURBO2 / 4) // vec FA iterations per block +typedef struct { + ggml_half norm; // 2 bytes: corrected L2 norm + uint8_t qs[QK_TURBO2 / 4]; // 8 bytes: 2-bit indices (4 per byte) +} block_turbo2_0; // 10 bytes total +static_assert(sizeof(block_turbo2_0) == sizeof(ggml_half) + QK_TURBO2/4, "wrong turbo2_0 block size/padding"); + +// opencoti-hook: turboquant tier family (M6-S2) — TurboQuant 8-bit KV, near-lossless rung. +// WHT-rotated, uniform signed int8 over the post-WHT normalized values (SCALE=256 ⇒ range +// ±0.496, step 1/256). 1 byte/value, no QJL signs, no codebook. block = group = 128. +// Per block: norm(fp16) + int8 indices (128 bytes) = 130 bytes per 128 values = 8.125 bits/value. +#define QK_TURBO8 128 +#define QK_TURBO8_GROUP 128 // rotation group size = head_dim +#define QR_TURBO8 1 // float2 dequant, same convention as the other turbo tiers +#define TURBO8_SCALE 256.0f // int8 quant scale: q = clamp(round(x*256), -127, 127) +typedef struct { + ggml_half norm; // 2 bytes: corrected L2 norm + int8_t qs[QK_TURBO8]; // 128 bytes: signed int8 (1 per value) +} block_turbo8_0; // 130 bytes total +static_assert(sizeof(block_turbo8_0) == sizeof(ggml_half) + QK_TURBO8, "wrong turbo8_0 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 +// = 4.0 bits/value +#define QK_TQ3_0 32 +typedef struct { + ggml_half d0; // 2 bytes: scale for first 16 elements + ggml_half d1; // 2 bytes: scale for last 16 elements + uint8_t qs[QK_TQ3_0 * 3 / 8]; // 12 bytes: 3-bit indices packed (4 groups of 8 in 3 bytes) +} block_tq3_1s; // 16 bytes total +static_assert(sizeof(block_tq3_1s) == 16, "wrong tq3_1s block size"); + +// TQ4_1S: WHT-rotated 4-bit weight quantization (16-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) + 4-bit indices packed (16 bytes) = 20 bytes per 32 values +// = 5.0 bits/value +#define QK_TQ4_1S 32 +typedef struct { + ggml_half d0; // 2 bytes: scale for first 16 elements + ggml_half d1; // 2 bytes: scale for last 16 elements + uint8_t qs[QK_TQ4_1S / 2]; // 16 bytes: 4-bit indices nibble-packed +} block_tq4_1s; // 20 bytes total +static_assert(sizeof(block_tq4_1s) == 20, "wrong tq4_1s block size"); + + #endif // GGML_COMMON_DECL #endif // GGML_COMMON_DECL diff --git a/llama.cpp/ggml/src/ggml-cpu/ggml-cpu.c b/llama.cpp/ggml/src/ggml-cpu/ggml-cpu.c --- a/llama.cpp/ggml/src/ggml-cpu/ggml-cpu.c +++ b/llama.cpp/ggml/src/ggml-cpu/ggml-cpu.c @@ -412,6 +412,27 @@ static const struct ggml_type_traits_cpu type_traits_cpu[GGML_TYPE_COUNT] = { [GGML_TYPE_I32] = { .from_float = (ggml_from_float_t) ggml_cpu_fp32_to_i32, }, + // opencoti-hook: turboquant (M6-S2) — turbo3 KV tail write needs a CPU from_float. In + // POSITION_WINDOW mode the host-pinned KV tail set_rows runs on the CPU backend; + // ggml_compute_forward_set_rows (ggml-cpu/ops.cpp) calls type_traits_cpu[dst->type].from_float, + // which was NULL for turbo3 → call 0x0 → PC=0 SIGSEGV at warmup (512-tok ubatch overflows the + // 256-cell window). quantize_row_turbo3_0_ref is the same forward-WHT+3bit+corrected-norm as the + // GPU set_rows kernel, so window/tail stay format-consistent for the dequant-on-lift. See poly_kv.md. + [GGML_TYPE_TURBO3_0] = { + .from_float = (ggml_from_float_t) quantize_row_turbo3_0_ref, + }, + // opencoti-hook: turboquant tier family (M6-S2) — same CPU from_float requirement as turbo3 + // for POSITION_WINDOW host-pinned tail set_rows. forward-WHT + Nbit + corrected-norm matches + // the GPU k_set_rows_turboN kernels ⇒ window/tail format-consistent for the dequant-on-lift. + [GGML_TYPE_TURBO4_0] = { + .from_float = (ggml_from_float_t) quantize_row_turbo4_0_ref, + }, + [GGML_TYPE_TURBO2_0] = { + .from_float = (ggml_from_float_t) quantize_row_turbo2_0_ref, + }, + [GGML_TYPE_TURBO8_0] = { + .from_float = (ggml_from_float_t) quantize_row_turbo8_0_ref, + }, }; const struct ggml_type_traits_cpu * ggml_get_type_traits_cpu(enum ggml_type type) { @@ -1753,6 +1774,16 @@ static void ggml_compute_forward(struct ggml_compute_params * params, struct ggm return; } + // opencoti diag (M6-S2): env-gated per-op trace to localize the turbo3 CPU crash; remove before ship. + if (params->ith == 0 && getenv("OPENCOTI_OPTRACE")) { + fprintf(stderr, "[OPTRACE] op=%s name=%s dst=%s s0=%s s1=%s\n", + ggml_op_name(tensor->op), tensor->name, + ggml_type_name(tensor->type), + tensor->src[0] ? ggml_type_name(tensor->src[0]->type) : "-", + tensor->src[1] ? ggml_type_name(tensor->src[1]->type) : "-"); + fflush(stderr); + } + // extra_buffer op? if (ggml_cpu_extra_compute_forward(params, tensor)) { return; @@ -2089,6 +2120,15 @@ static void ggml_compute_forward(struct ggml_compute_params * params, struct ggm // engine), a different path entirely. GGML_ABORT("STREAMING_FLASH_ATTN has no CPU implementation (CUDA-only op)"); } break; + case GGML_OP_TURBO_WHT: + { + // opencoti-hook: turboquant perf-lift (Level B) — F5 M6-S2 #391. + // CUDA-only op. The build_attn_mha fused-turbo gate emits it only on + // the GPU-resident turbo FA-VEC path (head_dim ≤ 256, K/V resident on + // CUDA), so — like STREAMING_FLASH_ATTN — the CPU scheduler never + // reaches here. supports_op keeps it on the CUDA backend. + GGML_ABORT("TURBO_WHT has no CPU implementation (CUDA-only op)"); + } break; case GGML_OP_GET_REL_POS: { ggml_compute_forward_get_rel_pos(params, tensor); @@ -2368,6 +2408,13 @@ static int ggml_get_n_tasks(struct ggml_tensor * node, int n_threads) { // scheduled on CPU. Single task keeps the planner well-formed. n_tasks = 1; } break; + case GGML_OP_TURBO_WHT: + { + // opencoti-hook: turboquant perf-lift (Level B) — F5 M6-S2 #391. + // CUDA-only; never scheduled on CPU. Single task keeps the planner + // well-formed. + n_tasks = 1; + } break; case GGML_OP_SILU_BACK: case GGML_OP_MUL: case GGML_OP_DIV: diff --git a/llama.cpp/ggml/src/ggml-cpu/quants.h b/llama.cpp/ggml/src/ggml-cpu/quants.h --- a/llama.cpp/ggml/src/ggml-cpu/quants.h +++ b/llama.cpp/ggml/src/ggml-cpu/quants.h @@ -19,6 +19,11 @@ void quantize_row_q5_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, in void quantize_row_q5_1(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); void quantize_row_q8_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); void quantize_row_q8_1(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +// opencoti-hook: turboquant (M6-S2) — turbo CPU from_float for POSITION_WINDOW host-tail KV write; see docs/features/poly_kv.md +void quantize_row_turbo3_0_ref(const float * GGML_RESTRICT x, block_turbo3_0 * GGML_RESTRICT y, int64_t k); +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); 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/convert.cu b/llama.cpp/ggml/src/ggml-cuda/convert.cu --- a/llama.cpp/ggml/src/ggml-cuda/convert.cu +++ b/llama.cpp/ggml/src/ggml-cuda/convert.cu @@ -1,5 +1,6 @@ #include "convert.cuh" #include "dequantize.cuh" +#include "turbo-dequant.cuh" // opencoti-hook: turboquant — see docs/features/poly_kv.md (M6-S2 Gate-2) #include @@ -726,6 +727,15 @@ to_fp16_cuda_t ggml_get_to_fp16_cuda(ggml_type type) { return dequantize_block_q8_0_f16_cuda; } return dequantize_block_cont_cuda; + // opencoti-hook: turboquant — see docs/features/poly_kv.md (M6-S2 Gate-2): inverse-WHT dequant → TRUE-space f16 + case GGML_TYPE_TURBO3_0: + return dequantize_block_cont_cuda; + case GGML_TYPE_TURBO4_0: + return dequantize_block_cont_cuda; + case GGML_TYPE_TURBO2_0: + return dequantize_block_cont_cuda; + case GGML_TYPE_TURBO8_0: + return dequantize_block_cont_cuda; case GGML_TYPE_Q2_K: return dequantize_row_q2_K_cuda; case GGML_TYPE_Q3_K: @@ -842,6 +852,15 @@ to_fp16_nc_cuda_t ggml_get_to_fp16_nc_cuda(ggml_type type) { return dequantize_block_cuda; case GGML_TYPE_Q8_0: return dequantize_block_cuda; + // opencoti-hook: turboquant — see docs/features/poly_kv.md (M6-S2 Gate-2): S3d lift path → TRUE-space f16 + case GGML_TYPE_TURBO3_0: + return dequantize_block_cuda; + case GGML_TYPE_TURBO4_0: + return dequantize_block_cuda; + case GGML_TYPE_TURBO2_0: + return dequantize_block_cuda; + case GGML_TYPE_TURBO8_0: + return dequantize_block_cuda; case GGML_TYPE_BF16: return convert_unary_cuda; default: diff --git a/llama.cpp/ggml/src/ggml-cuda/cpy.cu b/llama.cpp/ggml/src/ggml-cuda/cpy.cu --- a/llama.cpp/ggml/src/ggml-cuda/cpy.cu +++ b/llama.cpp/ggml/src/ggml-cuda/cpy.cu @@ -1,6 +1,8 @@ #include "cpy.cuh" #include "dequantize.cuh" #include "cpy-utils.cuh" +#include "turbo-cpy.cuh" // opencoti-hook: turboquant perf-lift (M6-S2 Level A) — GPU dequant-on-lift +#include "turbo-innerq-dev.cuh" // opencoti-hook: turboquant perf-lift (M6-S2 Level A) — InnerQ device scale accessor #if defined(GGML_USE_MUSA) && defined(GGML_MUSA_MUDNN_COPY) #include "ggml-musa/mudnn.cuh" #endif // GGML_USE_MUSA && GGML_MUSA_MUDNN_COPY @@ -558,6 +560,20 @@ void ggml_cuda_cpy(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, gg ggml_cpy_scalar_cuda (src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream); } + // opencoti-hook: turboquant perf-lift (M6-S2 Level A) — GPU-resident turbo→f16 dequant-on-lift. + // Replaces the CPU-spill cast (the scheduler used to assign ggml_cast(turboN→f16) to the CPU + // backend because CUDA CPY supports_op rejected turbo). The InnerQ device scale (owned by + // set-rows.cu) is fetched by the cast's row width (ne00 = n_embd_gqa); nullptr ⇒ identity + // (default-off / no slot). See docs/features/poly_kv.md (M6-S2 perf-path Level A). + } else if ((src0->type == GGML_TYPE_TURBO2_0 || src0->type == GGML_TYPE_TURBO3_0 || + src0->type == GGML_TYPE_TURBO4_0 || src0->type == GGML_TYPE_TURBO8_0) && + src1->type == GGML_TYPE_F16) { + const float * d_scale_inv = nullptr; + int head_dim = 0; + turbo_innerq_active_device_scale(ne00, &d_scale_inv, &head_dim); + ggml_cpy_turbo_f16_cuda(src0->type, src0_ddc, src1_ddc, ne, ne00, ne01, ne02, + nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, + d_scale_inv, head_dim, main_stream); } else { GGML_ABORT("%s: unsupported type combination (%s to %s)\n", __func__, ggml_type_name(src0->type), ggml_type_name(src1->type)); diff --git a/llama.cpp/ggml/src/ggml-cuda/fattn-common.cuh b/llama.cpp/ggml/src/ggml-cuda/fattn-common.cuh --- a/llama.cpp/ggml/src/ggml-cuda/fattn-common.cuh +++ b/llama.cpp/ggml/src/ggml-cuda/fattn-common.cuh @@ -584,10 +584,21 @@ static __device__ __forceinline__ void dequantize_V_q8_0(const void * __restrict } } +// opencoti-hook: turboquant perf-lift (Level B) — F5 M6-S2 #392. In-register turbo K/V readers +// (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" + template constexpr __device__ vec_dot_KQ_t get_vec_dot_KQ() { if constexpr (type_K == GGML_TYPE_F16) { return vec_dot_fattn_vec_KQ_f16; + } else if constexpr (type_K == GGML_TYPE_TURBO2_0) { // opencoti-hook: turboquant perf-lift (Level B) + return vec_dot_fattn_vec_KQ_turbo2; + } else if constexpr (type_K == GGML_TYPE_TURBO3_0) { + 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_Q4_0) { return vec_dot_fattn_vec_KQ_q4_0; } else if constexpr (type_K == GGML_TYPE_Q4_1) { @@ -610,6 +621,12 @@ template constexpr __device__ dequantize_V_t get_dequantize_V() { if constexpr (type_V == GGML_TYPE_F16) { return dequantize_V_f16; + } else if constexpr (type_V == GGML_TYPE_TURBO2_0) { // opencoti-hook: turboquant perf-lift (Level B) — F5 M6-S2 #392 + return dequantize_V_turbo2; + } else if constexpr (type_V == GGML_TYPE_TURBO3_0) { + return dequantize_V_turbo3; + } else if constexpr (type_V == GGML_TYPE_TURBO4_0) { + return dequantize_V_turbo4; } else if constexpr (type_V == GGML_TYPE_Q4_0) { return dequantize_V_q4_0; } else if constexpr (type_V == GGML_TYPE_Q4_1) { diff --git a/llama.cpp/ggml/src/ggml-cuda/fattn-turbo.cuh b/llama.cpp/ggml/src/ggml-cuda/fattn-turbo.cuh new file mode 100644 --- /dev/null +++ b/llama.cpp/ggml/src/ggml-cuda/fattn-turbo.cuh @@ -0,0 +1,129 @@ +// opencoti-hook: turboquant perf-lift (Level B) — see docs/features/poly_kv.md (F5 M6-S2 #392) +// +// In-register turbo K/V readers for the FUSED FA-VEC path. The turbo K/V cache stays compressed +// AND ROTATED on the GPU; these readers reconstruct centroid×norm (the WHT-rotated value) with NO +// inverse WHT — the Parseval identity makes Q_rot·K_rot == Q_true·K_true (Q is forward-rotated by +// the L3 ggml_turbo_wht(Q,0)), and the rotated attention output is un-rotated once by +// ggml_turbo_wht(O,1) at L3. InnerQ is therefore NOT applied here: it folds entirely into the Q +// pre-rotation (×scale_inv) and the output inverse-WHT (×scale_inv) — K and V reads are pure +// centroid×norm, exactly as the AtomicBot fork does. +// +// Turbo K is treated like F16 in our newer FA-VEC kernel: Q is kept FLOAT (Q_q8_1=false), and the +// turbo vec_dot mirrors vec_dot_fattn_vec_KQ_f16 — the per-half2 memcpy load becomes a centroid +// decode. A half2 pair (elements 2*e2, 2*e2+1) never straddles a 128-block (even start), and the +// V dequant i0 is 4-aligned, so block math is branch-free. Centroid tables are byte-identical to +// turbo-dequant.cuh / turbo-cpy.cu (validated in Level A). turbo8 / head_dim 512 are NOT here — +// they fall back to the Level-A materialize lift (fork has no turbo8, no D=512 vec instance). +#pragma once + +#include "common.cuh" + +// ---- centroid codebooks (== turbo-dequant.cuh TURBO{2,3,4}_CENTROIDS; N(0,1/128) Lloyd-Max) ---- +static __constant__ float FATTN_TURBO3_CENTROIDS[8] = { + -0.190685f, -0.117832f, -0.065717f, -0.021460f, 0.021460f, 0.065717f, 0.117832f, 0.190685f }; +static __constant__ float FATTN_TURBO4_CENTROIDS[16] = { + -0.173926f, -0.117195f, -0.089527f, -0.068756f, -0.051262f, -0.035597f, -0.020989f, -0.006938f, + 0.006938f, 0.020989f, 0.035597f, 0.051262f, 0.068756f, 0.089527f, 0.117195f, 0.173926f }; +static __constant__ float FATTN_TURBO2_CENTROIDS[4] = { + -0.133462f, -0.039994f, 0.039994f, 0.133462f }; + +// ---- per-element rotated-value decode (centroid[code]×norm). p = position within the 128-block. ---- +static __device__ __forceinline__ float fattn_turbo3_val(const block_turbo3_0 * blk, int p, float norm) { + const uint8_t lo = (blk->qs[p >> 2] >> ((p & 3) * 2)) & 0x3; + const uint8_t hi = (blk->signs[p >> 3] >> (p & 7)) & 0x1; + return FATTN_TURBO3_CENTROIDS[lo | (hi << 2)] * norm; +} +static __device__ __forceinline__ float fattn_turbo4_val(const block_turbo4_0 * blk, int p, float norm) { + const uint8_t idx = (blk->qs[p >> 1] >> ((p & 1) * 4)) & 0xF; + return FATTN_TURBO4_CENTROIDS[idx] * norm; +} +static __device__ __forceinline__ float fattn_turbo2_val(const block_turbo2_0 * blk, int p, float norm) { + const uint8_t idx = (blk->qs[p >> 2] >> ((p & 3) * 2)) & 0x3; + return FATTN_TURBO2_CENTROIDS[idx] * norm; +} + +// Q·K MAD against the float Q slice — Q_reg is half2 with V_DOT2_F32_F16_AVAILABLE, else float2 +// (mirrors vec_dot_fattn_vec_KQ_f16). kval is the decoded rotated-K half2 for this lane pair. +static __device__ __forceinline__ void fattn_turbo_qmad(float & sum, half2 kval, const void * Q_v, int qi) { +#ifdef V_DOT2_F32_F16_AVAILABLE + ggml_cuda_mad(sum, kval , ((const half2 *) Q_v)[qi]); +#else + ggml_cuda_mad(sum, __half22float2(kval), ((const float2 *) Q_v)[qi]); +#endif // V_DOT2_F32_F16_AVAILABLE +} + +// ===================== K: vec_dot (centroid×norm · float Q), f16-structured ===================== +// K_c → key row in turbo format; Q_v → this lane's float Q slice (Q_reg, already ×scale); the dot +// covers this thread's D-slice, warp_reduce_sum combines lanes (done by the caller). +#define FATTN_TURBO_KQ_BODY(BLOCK_T, VAL_FN) \ + 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; \ + _Pragma("unroll") \ + for (int k_KQ_0 = 0; k_KQ_0 < D/2; k_KQ_0 += nthreads*cpy_ne) { \ + half2 tmp[cpy_ne]; \ + _Pragma("unroll") \ + for (int k1 = 0; k1 < cpy_ne; ++k1) { \ + const int e2 = k_KQ_0 + (int)(threadIdx.x % nthreads)*cpy_ne + k1; \ + const int e = 2*e2; \ + const BLOCK_T * blk = (const BLOCK_T *) K_c + (e >> 7); \ + const float norm = __half2float(blk->norm); \ + const int ib = e & 127; \ + tmp[k1] = make_half2(VAL_FN(blk, ib, norm), VAL_FN(blk, ib + 1, norm)); \ + } \ + _Pragma("unroll") \ + for (int k1 = 0; k1 < cpy_ne; ++k1) { \ + fattn_turbo_qmad(sum, tmp[k1], Q_v, k_KQ_0/nthreads + k1); \ + } \ + } \ + return sum; + +template +static __device__ __forceinline__ float vec_dot_fattn_vec_KQ_turbo3( + const char * __restrict__ K_c, const void * __restrict__ Q_v, const int * __restrict__ Q_q8, const void * __restrict__ Q_ds_v) { + FATTN_TURBO_KQ_BODY(block_turbo3_0, fattn_turbo3_val) +} +template +static __device__ __forceinline__ float vec_dot_fattn_vec_KQ_turbo4( + const char * __restrict__ K_c, const void * __restrict__ Q_v, const int * __restrict__ Q_q8, const void * __restrict__ Q_ds_v) { + FATTN_TURBO_KQ_BODY(block_turbo4_0, fattn_turbo4_val) +} +template +static __device__ __forceinline__ float vec_dot_fattn_vec_KQ_turbo2( + const char * __restrict__ K_c, const void * __restrict__ Q_v, const int * __restrict__ Q_q8, const void * __restrict__ Q_ds_v) { + FATTN_TURBO_KQ_BODY(block_turbo2_0, fattn_turbo2_val) +} + +// ===================== V: dequant (centroid×norm), ne consecutive (i0 4-aligned) ===================== +// Reads ne consecutive rotated V values at element offset i0; writes T (half/float) as half2/float2. +#define FATTN_TURBO_V_BODY(BLOCK_T, VAL_FN) \ + static_assert(ne % 2 == 0, "bad ne"); \ + const BLOCK_T * blk = (const BLOCK_T *) vx + (i0 >> 7); \ + const float norm = __half2float(blk->norm); \ + const int ib = (int)(i0 & 127); \ + float vals[ne]; \ + _Pragma("unroll") \ + for (int l = 0; l < ne; ++l) vals[l] = VAL_FN(blk, ib + l, norm); \ + if constexpr (std::is_same_v) { \ + half2 * d2 = (half2 *) dst; \ + _Pragma("unroll") \ + for (int l = 0; l < ne; l += 2) d2[l/2] = make_half2(vals[l], vals[l+1]); \ + } else { \ + float2 * d2 = (float2 *) dst; \ + _Pragma("unroll") \ + for (int l = 0; l < ne; l += 2) d2[l/2] = make_float2(vals[l], vals[l+1]); \ + } + +template +static __device__ __forceinline__ void dequantize_V_turbo3(const void * __restrict__ vx, void * __restrict__ dst, const int64_t i0) { + FATTN_TURBO_V_BODY(block_turbo3_0, fattn_turbo3_val) +} +template +static __device__ __forceinline__ void dequantize_V_turbo4(const void * __restrict__ vx, void * __restrict__ dst, const int64_t i0) { + FATTN_TURBO_V_BODY(block_turbo4_0, fattn_turbo4_val) +} +template +static __device__ __forceinline__ void dequantize_V_turbo2(const void * __restrict__ vx, void * __restrict__ dst, const int64_t i0) { + FATTN_TURBO_V_BODY(block_turbo2_0, fattn_turbo2_val) +} diff --git a/llama.cpp/ggml/src/ggml-cuda/fattn-vec.cuh b/llama.cpp/ggml/src/ggml-cuda/fattn-vec.cuh --- a/llama.cpp/ggml/src/ggml-cuda/fattn-vec.cuh +++ b/llama.cpp/ggml/src/ggml-cuda/fattn-vec.cuh @@ -75,8 +75,12 @@ static __global__ void flash_attn_ext_vec( constexpr int nthreads_V_q = (D/4 < 32 ? D/4 : 32); #endif // GGML_USE_HIP + // 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 int nthreads = ggml_cuda_fattn_vec_get_nthreads_device(); - constexpr int nthreads_KQ = (type_K == GGML_TYPE_F16 || type_K == GGML_TYPE_BF16) ? 128 / cpy_nb : nthreads_KQ_q; + 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; static_assert(WARP_SIZE % nthreads_KQ == 0, "bad nthreads_K"); @@ -86,7 +90,9 @@ static __global__ void flash_attn_ext_vec( constexpr int V_cols_per_iter = WARP_SIZE / nthreads_V; constexpr vec_dot_KQ_t vec_dot_KQ = get_vec_dot_KQ(); - constexpr bool Q_q8_1 = type_K != GGML_TYPE_F16 && type_K != GGML_TYPE_BF16; + // opencoti-hook: turboquant perf-lift (Level B) — F5 M6-S2 #392. Keep Q FLOAT for turbo K + // (q8_1-quantizing Q would destroy the WHT-rotation precision the Parseval identity needs). + constexpr bool Q_q8_1 = type_K != GGML_TYPE_F16 && type_K != GGML_TYPE_BF16 && !is_turbo_K; #ifdef V_DOT2_F32_F16_AVAILABLE constexpr dequantize_V_t dequantize_V = get_dequantize_V(); #else @@ -601,3 +607,14 @@ EXTERN_DECL_FATTN_VEC_CASES(256, GGML_TYPE_Q5_0) 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) + +// opencoti-hook: turboquant perf-lift (Level B) — F5 M6-S2 #392. Fused turbo K/V vec instances +// (K==V same tier only; instances in template-instances/fattn-vec-instance-turbo{2,3,4}_0-*.cu). +#define EXTERN_DECL_FATTN_VEC_TURBO(type_T) \ + extern DECL_FATTN_VEC_CASE( 64, type_T, type_T); \ + extern DECL_FATTN_VEC_CASE(128, type_T, type_T); \ + extern DECL_FATTN_VEC_CASE(256, type_T, type_T); \ + +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) diff --git a/llama.cpp/ggml/src/ggml-cuda/fattn.cu b/llama.cpp/ggml/src/ggml-cuda/fattn.cu --- a/llama.cpp/ggml/src/ggml-cuda/fattn.cu +++ b/llama.cpp/ggml/src/ggml-cuda/fattn.cu @@ -330,6 +330,13 @@ 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: 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). + FATTN_VEC_CASES_ALL_D(GGML_TYPE_TURBO2_0, GGML_TYPE_TURBO2_0) + 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) + GGML_ABORT("fatal error"); } @@ -446,6 +453,10 @@ static best_fattn_kernel ggml_cuda_get_best_fattn_kernel(const int device, const case GGML_TYPE_Q8_0: 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: + break; default: return BEST_FATTN_KERNEL_NONE; } @@ -458,6 +469,14 @@ static best_fattn_kernel ggml_cuda_get_best_fattn_kernel(const int device, const // 192 satisfies % 64 == 0 but has no vec instance (DKQ != DV); force it onto the MMA path. const bool can_use_vector_kernel = Q->ne[0] <= 256 && Q->ne[0] % 64 == 0 && Q->ne[0] != 192 && K->ne[1] % FATTN_KQ_STRIDE == 0; + // opencoti-hook: turboquant perf-lift (Level B) — F5 M6-S2 #392. Turbo K/V have ONLY a fused + // 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) { + return can_use_vector_kernel ? BEST_FATTN_KERNEL_VEC : BEST_FATTN_KERNEL_NONE; + } + // If Turing tensor cores are available, use them: if (turing_mma_available(cc) && Q->ne[0] != 40 && Q->ne[0] != 72) { if (can_use_vector_kernel) { diff --git a/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu b/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu --- a/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu @@ -68,6 +68,7 @@ #include "ggml-cuda/tri.cuh" #include "ggml-cuda/cumsum.cuh" #include "ggml-cuda/fill.cuh" +#include "ggml-cuda/turbo-wht.cuh" // opencoti-hook: turboquant perf-lift (Level B) — F5 M6-S2 #391 #include "ggml.h" #include @@ -3039,6 +3040,13 @@ static bool ggml_cuda_compute_forward(ggml_backend_cuda_context & ctx, struct gg // here in F32, not delegated to the kernel. See fattn.cu. ggml_cuda_streaming_flash_attn(ctx, dst); break; + case GGML_OP_TURBO_WHT: + // opencoti-hook: turboquant perf-lift (Level B) — F5 M6-S2 #391. + // Walsh–Hadamard rotation of the small Q (fwd) / FA-output (inv) tensors + // for the fused turbo FA-VEC path. InnerQ resolved device-side. See + // turbo-wht.cu. + ggml_cuda_turbo_wht(ctx, dst); + break; case GGML_OP_OUT_PROD: ggml_cuda_out_prod(ctx, dst); break; @@ -5375,9 +5383,13 @@ static bool GGML_CALL ggml_backend_cuda_device_supports_op(ggml_backend_dev_t de } break; case GGML_OP_SET_ROWS: { + // opencoti-hook: turboquant tier family — see docs/features/poly_kv.md (M6-S2): + // accept turbo2/turbo3/turbo4 dst (each has a GPU k_set_rows_turboN write kernel). 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) && + op->type == GGML_TYPE_Q5_1 || op->type == GGML_TYPE_Q8_0 || op->type == GGML_TYPE_IQ4_NL || + 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->src[0]->type == GGML_TYPE_F32 && (op->src[1]->type == GGML_TYPE_I64 || op->src[1]->type == GGML_TYPE_I32); } break; @@ -5397,6 +5409,15 @@ static bool GGML_CALL ggml_backend_cuda_device_supports_op(ggml_backend_dev_t de ) { return true; } + // opencoti-hook: turboquant perf-lift (M6-S2 Level A) — GPU-resident turbo→f16 + // dequant-on-lift (ggml_cuda_cpy → ggml_cpy_turbo_f16_cuda). Lets the scheduler keep + // build_attn_mha's ggml_cast(turboN→f16) on CUDA instead of spilling to the CPU + // backend. See docs/features/poly_kv.md (M6-S2 perf-path Level A). + if ((src0_type == GGML_TYPE_TURBO2_0 || src0_type == GGML_TYPE_TURBO3_0 || + src0_type == GGML_TYPE_TURBO4_0 || src0_type == GGML_TYPE_TURBO8_0) && + src1_type == GGML_TYPE_F16) { + return true; + } if (src0_type == GGML_TYPE_F32 && src1_type == GGML_TYPE_Q8_0) { return true; } @@ -5623,6 +5644,18 @@ static bool GGML_CALL ggml_backend_cuda_device_supports_op(ggml_backend_dev_t de return false; } } + case GGML_OP_TURBO_WHT: { + // opencoti-hook: turboquant perf-lift (Level B) — F5 M6-S2 #391. f32 in/out, + // contiguous, head_dim (ne[0]) a multiple of the 128 WHT group. group_size is + // stored in op_params[1] (the ctor forces 128 for our path). InnerQ resolved + // device-side; no src[1] requirement. + const ggml_tensor * a = op->src[0]; + if (a->type != GGML_TYPE_F32 || op->type != GGML_TYPE_F32) return false; + if (!ggml_is_contiguous(a) || !ggml_is_contiguous(op)) return false; + int group_size = 0; + memcpy(&group_size, (const char *) op->op_params + sizeof(int), sizeof(int)); + return group_size == 128 && a->ne[0] % 128 == 0; + } case GGML_OP_CROSS_ENTROPY_LOSS: case GGML_OP_CROSS_ENTROPY_LOSS_BACK: case GGML_OP_OPT_STEP_ADAMW: diff --git a/llama.cpp/ggml/src/ggml-cuda/set-rows.cu b/llama.cpp/ggml/src/ggml-cuda/set-rows.cu --- a/llama.cpp/ggml/src/ggml-cuda/set-rows.cu +++ b/llama.cpp/ggml/src/ggml-cuda/set-rows.cu @@ -1,5 +1,10 @@ #include "set-rows.cuh" #include "cpy-utils.cuh" +#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 +#include typedef void (*set_rows_kernel_t)(const char * src, char * dst); @@ -214,6 +219,801 @@ static void set_rows_cuda( } } +// opencoti-hook: turboquant — see docs/features/poly_kv.md (M6-S2 Gate-2 write path). +// f32 K/V → turbo3: group-level forward WHT + 3-bit centroid quantize + corrected norm. +// One block per 128-element WHT group (QK_TURBO3 == group). InnerQ stub-disabled +// (Gate-3 re-enables). No tail kernel — head_dim must be 128-block-aligned (Qwen 128, +// Gemma 256 = 2x128). Algorithm + warp/ballot packing ported from the AtomicBot fork's +// ggml-cuda/set-rows.cu k_set_rows_turbo3, with the InnerQ calibrate/apply blocks removed. +// ===== opencoti-hook: turboquant InnerQ (M6-S2) — per-channel equalization ===== +// Device state — set-rows.cu is the ONLY TU with device access (the Qwen gate uses +// the host CPU dequant; cross-TU GPU dequant InnerQ for the iSWA streaming path is a +// deferred follow-up). __device__ globals are zero-initialized ⇒ OFF mode is inert. +// opencoti-hook: turboquant InnerQ (M6-S2) — PER-LAYER (per-width) head_dim folding. +// ne00 of the KV set-rows row is n_embd_gqa = head_dim*n_head_kv (NOT head_dim). A model +// can mix head_dims across layers: Gemma-4 = 256-SWA (n_embd_gqa 256*8=2048) + 512-global +// (512*2=1024). head_dim is NOT derivable from ne00 alone (ambiguous), and is unavailable +// at the CUDA launch site, so the host supplies a {n_embd_gqa→head_dim} map via env +// (TURBO_INNERQ_HEADDIM_MAP="2048:256,1024:512"). Each distinct width is a SLOT with its +// own scale_inv[head_dim] and sq_accum — no cross-width contamination. The dispatch maps +// ne00→slot and passes it as a kernel arg (no per-launch H2D). The per-head channel is +// (i_grp*128 + j) % head_dim[slot]; scale shared across heads of that width. head_dim==128 +// ⇒ fold is identity (j). Empty map ⇒ one wildcard slot (back-compat, uniform models). +static __device__ float d_innerq_scale [INNERQ_MAX_SLOTS][INNERQ_MAX_CHANNELS]; +static __device__ float d_innerq_scale_inv[INNERQ_MAX_SLOTS][INNERQ_MAX_CHANNELS]; +static __device__ float d_innerq_sq_accum [INNERQ_MAX_SLOTS][INNERQ_MAX_CHANNELS]; +static __device__ int d_innerq_slot_hd [INNERQ_MAX_SLOTS]; // head_dim per slot +static __device__ int d_innerq_count; // total groups seen (finalize trigger) +static __device__ int d_innerq_active; // 0 = identity, 1 = scales applied +static __device__ int d_innerq_calibrating; // 1 = accumulating K² stats + +static int innerq_mode = 0; // 0=off, 1=calibrating, 2=active +static int innerq_target_blocks = 0; +static float innerq_strength = 0.5f; +static bool innerq_initialized = false; +static const char * innerq_out_path = nullptr; +// Per-width slot table (host mirror of the device map). slot_nembd[s]==0 ⇒ wildcard. +static int innerq_nslots = 0; +static int innerq_slot_nembd[INNERQ_MAX_SLOTS] = {0}; +static int innerq_slot_hd [INNERQ_MAX_SLOTS] = {0}; + +// Map a set-rows row width (ne00 = n_embd_gqa) to its InnerQ slot, or -1 (no InnerQ). +// Exact width match first, then a wildcard (nembd==0) slot for uniform/back-compat. +static int innerq_slot_for(int ne00) { + for (int s = 0; s < innerq_nslots; s++) if (innerq_slot_nembd[s] == ne00) return s; + for (int s = 0; s < innerq_nslots; s++) if (innerq_slot_nembd[s] == 0) return s; + return -1; +} + +// Parse TURBO_INNERQ_HEADDIM_MAP="ne0:hd,ne1:hd,...". Returns nslots (0 if unset/empty). +// On empty/parse-fail the caller falls back to a single wildcard slot from HEAD_DIM. +static int innerq_parse_map(int * nembd, int * hd) { + const char * m = getenv("TURBO_INNERQ_HEADDIM_MAP"); + if (!m || !m[0]) return 0; + int n = 0; + const char * p = m; + while (*p && n < INNERQ_MAX_SLOTS) { + char * end = nullptr; + long ne = strtol(p, &end, 10); + if (end == p || *end != ':') break; + p = end + 1; + long h = strtol(p, &end, 10); + if (end == p) break; + if (ne > 0 && h >= 128 && h <= INNERQ_MAX_CHANNELS && (h % 128) == 0) { + nembd[n] = (int) ne; hd[n] = (int) h; n++; + } + p = end; + while (*p == ',' || *p == ' ') p++; + } + return n; +} + +// ACTIVE-from-token-0: upload per-slot scale (=1/scale_inv) + scale_inv + head_dim, +// publish the host slot table, activate. scale_inv_in[s][0..hd[s]). +static void turbo_innerq_upload_active(int nslots, const int * nembd, const int * hd, + const float scale_inv_in[][INNERQ_MAX_CHANNELS]) { + float scale[INNERQ_MAX_SLOTS][INNERQ_MAX_CHANNELS]; + float sinv [INNERQ_MAX_SLOTS][INNERQ_MAX_CHANNELS]; + int hd_dev[INNERQ_MAX_SLOTS]; + for (int s = 0; s < INNERQ_MAX_SLOTS; s++) { + const int h = (s < nslots) ? hd[s] : 0; + hd_dev[s] = h; + for (int i = 0; i < INNERQ_MAX_CHANNELS; i++) { + const float si = (s < nslots && i < h) ? scale_inv_in[s][i] : 1.0f; + sinv [s][i] = si; + scale[s][i] = (si > 1e-20f) ? (1.0f / si) : 1.0f; + } + } + // publish host mirror (slot selection at dispatch) + innerq_nslots = nslots; + for (int s = 0; s < INNERQ_MAX_SLOTS; s++) { + innerq_slot_nembd[s] = (s < nslots) ? nembd[s] : 0; + innerq_slot_hd[s] = (s < nslots) ? hd[s] : 0; + } + int zero = 0, one = 1; + CUDA_CHECK(cudaMemcpyToSymbol(d_innerq_calibrating, &zero, sizeof(int))); + CUDA_CHECK(cudaMemcpyToSymbol(d_innerq_scale, scale, sizeof(scale))); + CUDA_CHECK(cudaMemcpyToSymbol(d_innerq_scale_inv, sinv, sizeof(sinv))); + CUDA_CHECK(cudaMemcpyToSymbol(d_innerq_slot_hd, hd_dev, sizeof(hd_dev))); + CUDA_CHECK(cudaDeviceSynchronize()); + CUDA_CHECK(cudaMemcpyToSymbol(d_innerq_active, &one, sizeof(int))); +} + +static void turbo_innerq_init(void) { + if (innerq_initialized) return; + innerq_initialized = true; + + const char * scale_path = getenv("TURBO_INNERQ_SCALE"); + if (scale_path && scale_path[0]) { + static float sinv[INNERQ_MAX_SLOTS][INNERQ_MAX_CHANNELS]; + int nembd[INNERQ_MAX_SLOTS], hd[INNERQ_MAX_SLOTS]; + const int ns = turbo_innerq_load_file_v2(scale_path, nembd, hd, sinv); + if (ns > 0) { + turbo_innerq_upload_active(ns, nembd, hd, sinv); + innerq_mode = 2; + GGML_LOG_INFO("turbo_innerq_init: InnerQ ACTIVE from token 0 (%d slot(s) from %s)\n", ns, scale_path); + for (int s = 0; s < ns; s++) + GGML_LOG_INFO("turbo_innerq_init: slot %d: n_embd_gqa=%d head_dim=%d\n", s, nembd[s], hd[s]); + } else { + GGML_LOG_WARN("turbo_innerq_init: InnerQ scale file %s unreadable — disabled\n", scale_path); + } + return; + } + + const char * cal = getenv("TURBO_INNERQ_CALIBRATE"); + if (cal && atoi(cal) > 0) { + innerq_target_blocks = atoi(cal); + innerq_out_path = getenv("TURBO_INNERQ_OUT"); + const char * str = getenv("TURBO_INNERQ_STRENGTH"); + if (str) innerq_strength = (float) atof(str); + if (innerq_strength <= 0.0f || innerq_strength > 1.0f) innerq_strength = 0.5f; + // Per-width slots: TURBO_INNERQ_HEADDIM_MAP="2048:256,1024:512" (Gemma mixed). For + // uniform models leave it unset and pass TURBO_INNERQ_HEAD_DIM (default 128) ⇒ one + // wildcard slot (n_embd_gqa=0 matches any width). Each slot folds per its head_dim. + innerq_nslots = innerq_parse_map(innerq_slot_nembd, innerq_slot_hd); + if (innerq_nslots <= 0) { + const char * hd_env = getenv("TURBO_INNERQ_HEAD_DIM"); + int hd = (hd_env && hd_env[0]) ? atoi(hd_env) : 128; + if (hd < 128 || hd > INNERQ_MAX_CHANNELS || (hd % 128) != 0) hd = 128; + innerq_nslots = 1; innerq_slot_nembd[0] = 0; innerq_slot_hd[0] = hd; + } + int slot_hd_dev[INNERQ_MAX_SLOTS] = {0}; + for (int s = 0; s < innerq_nslots; s++) slot_hd_dev[s] = innerq_slot_hd[s]; + float zeros[INNERQ_MAX_SLOTS][INNERQ_MAX_CHANNELS] = {{0}}; + int zero = 0, one = 1; + CUDA_CHECK(cudaMemcpyToSymbol(d_innerq_sq_accum, zeros, sizeof(zeros))); + CUDA_CHECK(cudaMemcpyToSymbol(d_innerq_count, &zero, sizeof(int))); + CUDA_CHECK(cudaMemcpyToSymbol(d_innerq_active, &zero, sizeof(int))); + CUDA_CHECK(cudaMemcpyToSymbol(d_innerq_slot_hd, slot_hd_dev, sizeof(slot_hd_dev))); + CUDA_CHECK(cudaMemcpyToSymbol(d_innerq_calibrating, &one, sizeof(int))); + innerq_mode = 1; + GGML_LOG_INFO("turbo_innerq_init: InnerQ CALIBRATING (target=%d blocks, strength=%.2f, %d slot(s), out=%s)\n", + innerq_target_blocks, innerq_strength, innerq_nslots, innerq_out_path ? innerq_out_path : "(none)"); + for (int s = 0; s < innerq_nslots; s++) + GGML_LOG_INFO("turbo_innerq_init: slot %d: n_embd_gqa=%d head_dim=%d\n", s, innerq_slot_nembd[s], innerq_slot_hd[s]); + } +} + +// opencoti-hook: turboquant perf-lift (M6-S2 Level A) — cross-TU device-scale accessor. +// Exposes the active slot's scale_inv device row + head_dim to the GPU dequant-on-lift cpy kernel +// (turbo-cpy.cu), selecting the slot by the cast's row width (ne00 = n_embd_gqa) exactly as the CPU +// dequant does (turbo_innerq_cpu_slot_for). Returns false when InnerQ is not ACTIVE or no slot +// matches ⇒ the kernel un-equalizes with identity (no ×scale_inv), matching the CPU miss path. +bool turbo_innerq_active_device_scale(int64_t ne00, const float ** d_scale_inv_out, int * head_dim_out) { + turbo_innerq_init(); // idempotent (innerq_initialized guard) + if (innerq_mode != 2) return false; // 2 = ACTIVE (not off/calibrating) + const int slot = innerq_slot_for((int) ne00); + if (slot < 0) return false; + void * base = nullptr; + if (cudaGetSymbolAddress(&base, d_innerq_scale_inv) != cudaSuccess || base == nullptr) return false; + *d_scale_inv_out = (const float *) base + (size_t) slot * INNERQ_MAX_CHANNELS; + *head_dim_out = innerq_slot_hd[slot]; + return true; +} + +static void turbo_innerq_finalize(int group_size) { + // Per-slot finalize. Within a slot, every head-channel is written the same number of + // times (n_head_kv*n_tokens*n_layers_of_that_width), so that uniform count cancels in + // the mean_rms/rms ratio — no per-slot count normalization needed (a single global + // `count` is only the finalize trigger). Slots are independent (no cross-width mixing). + (void) group_size; + float sq_accum[INNERQ_MAX_SLOTS][INNERQ_MAX_CHANNELS]; + int count = 0; + CUDA_CHECK(cudaMemcpyFromSymbol(sq_accum, d_innerq_sq_accum, sizeof(sq_accum))); + CUDA_CHECK(cudaMemcpyFromSymbol(&count, d_innerq_count, sizeof(int))); + int zero = 0; + CUDA_CHECK(cudaMemcpyToSymbol(d_innerq_calibrating, &zero, sizeof(int))); // stop accumulating + innerq_mode = 0; + if (count <= 0) { + GGML_LOG_WARN("turbo_innerq_finalize: 0 blocks accumulated — nothing written\n"); + return; + } + // opencoti-hook: turboquant InnerQ (M6-S2) — env-tunable equalization clamp. + // The original hard [0.5,2.0] cap (max 2x) under-equalizes models with extreme + // per-channel K outliers: Qwen2.5-1.5B has ~10/128 K channels needing >=4x (ratio>=4 + // ⇒ scale>=2), which the 2x cap saturated, leaving turbo3-K too lossy for softmax + // while turbo3-V (smoothing-tolerant) stayed fine. Widen to a 4x default and allow + // TURBO_INNERQ_CLAMP=C (scale ∈ [1/C, C]) to sweep without a rebuild (re-read at each + // calibration finalize). C clamps total equalization swing; strength shapes its curve. + float clamp_hi = 4.0f; + { + const char * ce = getenv("TURBO_INNERQ_CLAMP"); + if (ce && ce[0]) { const float v = (float) atof(ce); if (v >= 1.0f) clamp_hi = v; } + } + const float clamp_lo = 1.0f / clamp_hi; + static float scale_inv[INNERQ_MAX_SLOTS][INNERQ_MAX_CHANNELS]; + int out_nembd[INNERQ_MAX_SLOTS], out_hd[INNERQ_MAX_SLOTS]; + for (int s = 0; s < innerq_nslots; s++) { + const int n_ch = innerq_slot_hd[s]; + out_nembd[s] = innerq_slot_nembd[s]; + out_hd[s] = n_ch; + float rms[INNERQ_MAX_CHANNELS], mean_rms = 0.0f, max_ratio = 0.0f, min_ratio = 1e30f; + for (int i = 0; i < n_ch; i++) { rms[i] = sqrtf(sq_accum[s][i] / (float)count); mean_rms += rms[i]; } + mean_rms /= (float)n_ch; + int n_clamped_hi = 0; + for (int i = 0; i < n_ch; i++) { + const float ratio = (rms[i] > 1e-10f) ? (mean_rms / rms[i]) : 1.0f; + float sc = powf(ratio, innerq_strength); + if (sc < clamp_lo) sc = clamp_lo; + if (sc > clamp_hi) { sc = clamp_hi; n_clamped_hi++; } + scale_inv[s][i] = 1.0f / sc; + if (ratio > max_ratio) max_ratio = ratio; + if (ratio < min_ratio) min_ratio = ratio; + } + GGML_LOG_INFO("turbo_innerq_finalize: slot %d (n_embd_gqa=%d head_dim=%d) clamp=[%.3f,%.3f] strength=%.2f → " + "%d/%d saturated, max_ratio=%.3f min_ratio=%.3f\n", + s, out_nembd[s], n_ch, clamp_lo, clamp_hi, innerq_strength, n_clamped_hi, n_ch, max_ratio, min_ratio); + } + if (innerq_out_path && innerq_out_path[0]) { + const bool ok = turbo_innerq_dump_file_v2(innerq_out_path, innerq_nslots, out_nembd, out_hd, scale_inv); + GGML_LOG_INFO("turbo_innerq_finalize: calibrated on %d blocks, %d slot(s) → %s %s\n", + count, innerq_nslots, innerq_out_path, ok ? "OK" : "WRITE-FAILED"); + } else { + GGML_LOG_WARN("turbo_innerq_finalize: calibrated but TURBO_INNERQ_OUT unset — nothing written\n"); + } +} + +// Called before each turbo3 set-rows launch (host). Lazily inits; finalizes when +// the calibration block target is reached. +static void turbo_innerq_check_finalize(int group_size) { + if (!innerq_initialized) turbo_innerq_init(); + if (innerq_mode != 1) return; // only calibrating finalizes + if (group_size < 128) { // InnerQ math assumes WHT group == head_dim + GGML_LOG_WARN("turbo_innerq_check_finalize: disabled (group_size=%d < 128)\n", group_size); + innerq_mode = 0; int zero = 0; + CUDA_CHECK(cudaMemcpyToSymbol(d_innerq_calibrating, &zero, sizeof(int))); + return; + } + int count = 0; + CUDA_CHECK(cudaMemcpyFromSymbol(&count, d_innerq_count, sizeof(int))); + if (count >= innerq_target_blocks) turbo_innerq_finalize(group_size); +} +// ===== end InnerQ ===== + +template +__launch_bounds__(128) +static __global__ void k_set_rows_turbo3( + const float * __restrict__ src0, + const idx_t * __restrict__ src1, + block_turbo3_0 * __restrict__ dst, + const int64_t ne00, const int64_t ne01, const int64_t ne02, + const int64_t ne11, const int64_t ne12, + 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 int64_t s1, const int64_t s2, const int64_t s3, + const int innerq_slot) { // InnerQ per-width slot (-1 = off) + constexpr int GROUP_SIZE = QK_TURBO3; // 128 + const int j = threadIdx.x; // 0..127, element within group + + // Decode group index → (i_grp, i01, i02, i03). MUST match this tree's + // k_set_rows_quant convention (i02 % ne02, i12 = i03 % ne12, i11 = i02 % ne11), + // NOT the AtomicBot fork's (which assumes ne02==ne12). + const int64_t n_groups_per_row = ne00 / GROUP_SIZE; + const int64_t g = blockIdx.x; + const int64_t i_grp = g % n_groups_per_row; + int64_t tmp = g / n_groups_per_row; + const int64_t i01 = tmp % ne01; + tmp = tmp / ne01; + const int64_t i02 = tmp % ne02; + const int64_t i03 = tmp / ne02; + + const int64_t i12 = i03 % ne12; + const int64_t i11 = i02 % ne11; + const int64_t i10 = i01; + + const int64_t dst_row = *(src1 + i10*s10 + i11*s11 + i12*s12); + const float * src_row = src0 + i01*s01 + i02*s02 + i03*s03; + block_turbo3_0 * dst_row_ptr = (block_turbo3_0 *)((char *)dst + dst_row*s1 + i02*s2 + i03*s3); + block_turbo3_0 * blk = dst_row_ptr + i_grp; // blocks_per_group == 1 + + __shared__ float x[GROUP_SIZE]; + x[j] = src_row[i_grp * GROUP_SIZE + j]; + __syncthreads(); + + // opencoti-hook: turboquant InnerQ (M6-S2) — calibrate on RAW values, then equalize. + // Calibrate (accumulate K² per channel) and apply (×scale) are mutually exclusive + // (calibrate-then-activate flow): calibration runs un-equalized, active runs equalize + // from token 0. Both gated on zero-initialized device flags ⇒ OFF mode is inert. + // head_dim-aware, PER-WIDTH SLOT: the per-head channel is (i_grp*128 + j) % head_dim + // of this layer's width-slot (scale shared across heads of that width). innerq_slot is + // resolved host-side from ne00; -1 ⇒ no InnerQ for this width. head_dim==128 ⇒ ch==j + // (identity, byte-compatible with the old uniform path). + const int innerq_hd = (innerq_slot >= 0) ? d_innerq_slot_hd[innerq_slot] : 0; + if (innerq_slot >= 0 && d_innerq_calibrating && innerq_hd > 0) { + const int ch = (int)((i_grp * GROUP_SIZE + j) % innerq_hd); + atomicAdd(&d_innerq_sq_accum[innerq_slot][ch], x[j] * x[j]); + if (j == 0) atomicAdd(&d_innerq_count, 1); + } + if (innerq_slot >= 0 && d_innerq_active && innerq_hd > 0) { + const int ch = (int)((i_grp * GROUP_SIZE + j) % innerq_hd); + x[j] *= d_innerq_scale[innerq_slot][ch]; + } + __syncthreads(); + + // Step 2: parallel L2 norm (warp reduce + inter-warp via shared) + constexpr int n_warps = GROUP_SIZE / WARP_SIZE; + __shared__ float warp_accum[n_warps]; + float v2 = x[j] * x[j]; + for (int offset = WARP_SIZE/2; offset > 0; offset >>= 1) v2 += __shfl_xor_sync(0xffffffff, v2, offset); + if (j % WARP_SIZE == 0) warp_accum[j / WARP_SIZE] = v2; + __syncthreads(); + __shared__ float s_norm_sq; + if (j == 0) { float t = 0.0f; for (int w = 0; w < n_warps; w++) t += warp_accum[w]; s_norm_sq = t; } + __syncthreads(); + const float grp_norm = sqrtf(s_norm_sq); + const float inv_norm = (grp_norm > 1e-10f) ? 1.0f / grp_norm : 0.0f; + + // Step 3: normalize + x[j] *= inv_norm; + __syncthreads(); + + // Step 4: forward WHT (signs1 → butterfly → 1/sqrt128 · signs2) + x[j] *= TURBO3_WHT_SIGNS1[j]; + __syncthreads(); +#define WHT_STAGE_SHARED(h) \ + if (j % (2*(h)) < (h)) { float a = x[j], b = x[j+(h)]; x[j] = a+b; x[j+(h)] = a-b; } \ + __syncthreads(); + WHT_STAGE_SHARED(1) WHT_STAGE_SHARED(2) WHT_STAGE_SHARED(4) + WHT_STAGE_SHARED(8) WHT_STAGE_SHARED(16) WHT_STAGE_SHARED(32) WHT_STAGE_SHARED(64) +#undef WHT_STAGE_SHARED + const float inv_sqrt_group = 0.08838834764831845f; // 1/sqrt(128) + x[j] = x[j] * inv_sqrt_group * TURBO3_WHT_SIGNS2[j]; + __syncthreads(); + + // Step 5: quantize element j → 3-bit centroid index + const uint8_t idx = turbo3_nearest_centroid_3bit(x[j]); + + // Step 6: warp-cooperative pack (no atomics) + const int lane = j % WARP_SIZE; + const int elem_in_block = j; // QK_TURBO3 == GROUP_SIZE, single block + const int qs_byte_idx = elem_in_block / 4; + const uint8_t my_low2 = idx & 0x3; + uint8_t qs_byte = 0; +#pragma unroll + for (int k = 0; k < 4; k++) { + uint8_t contrib = __shfl_sync(0xffffffff, my_low2, (lane & ~3) + k); + qs_byte |= contrib << (k * 2); + } + if (lane % 4 == 0) blk->qs[qs_byte_idx] = qs_byte; + + const uint32_t ballot = __ballot_sync(0xffffffff, (idx >> 2) & 1); + const int local_signs_byte = lane / 8; + const int global_signs_byte = elem_in_block / 8; + const uint8_t signs_byte = (uint8_t)((ballot >> (local_signs_byte * 8)) & 0xFF); + if (lane % 8 == 0) blk->signs[global_signs_byte] = signs_byte; + + // Step 7: reconstruction norm (parallel, reuse warp_accum) + const float c = TURBO3_CENTROIDS_3BIT[idx]; + float rc = c * c; + for (int offset = WARP_SIZE/2; offset > 0; offset >>= 1) rc += __shfl_xor_sync(0xffffffff, rc, offset); + if (j % WARP_SIZE == 0) warp_accum[j / WARP_SIZE] = rc; + __syncthreads(); + __shared__ float s_recon_sq; + if (j == 0) { float t = 0.0f; for (int w = 0; w < n_warps; w++) t += warp_accum[w]; s_recon_sq = t; } + __syncthreads(); + const float recon_norm = sqrtf(s_recon_sq); + const float corrected_norm = (recon_norm > 1e-10f) ? grp_norm / recon_norm : grp_norm; + + // Step 8: write corrected norm (one thread per block) + if (elem_in_block == 0) blk->norm = __float2half(corrected_norm); +} + +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; + const idx_t * src1_d = (const idx_t *)src1->data; + + GGML_TENSOR_BINARY_OP_LOCALS + GGML_ASSERT(ne00 % QK_TURBO3 == 0); // 128-block aligned (Gate-2: no tail kernel) + + cudaStream_t stream = ctx.stream(); + const int64_t s01 = nb01 / sizeof(float), s02 = nb02 / sizeof(float), s03 = nb03 / sizeof(float); + const int64_t s10 = nb10 / sizeof(idx_t), s11 = nb11 / sizeof(idx_t), s12 = nb12 / sizeof(idx_t); + const int64_t n_groups = ne00 / QK_TURBO3; + const int64_t ne_total = n_groups * ne01 * ne02 * ne03; + if (ne_total <= 0) return; + + // opencoti-hook: turboquant InnerQ (M6-S2) — lazy init + finalize calibration, then + // resolve this layer's per-width slot from ne00 (= n_embd_gqa). -1 ⇒ no InnerQ here. + turbo_innerq_check_finalize(QK_TURBO3); + const int innerq_slot = (innerq_mode != 0) ? innerq_slot_for((int) ne00) : -1; + + k_set_rows_turbo3<<<(int)ne_total, QK_TURBO3, 0, stream>>>( + src0_d, src1_d, (block_turbo3_0 *)dst->data, + ne00, ne01, ne02, ne11, ne12, + s01, s02, s03, s10, s11, s12, + nb1, nb2, nb3, innerq_slot); +} + +// ===================== turbo4 (4-bit nibble) write kernel ===================== +// opencoti-hook: turboquant — see docs/features/poly_kv.md (M6-S2 tier family). Clone of +// k_set_rows_turbo3: identical InnerQ / L2-norm / forward-WHT preamble (width-independent), +// differs only at quantize (4-bit centroid), pack (nibble, 2/byte, no QJL signs) and the +// reconstruction-norm centroid table. block = group = QK_TURBO4 = 128. +template +static __global__ void k_set_rows_turbo4( + const float * __restrict__ src0, + const idx_t * __restrict__ src1, + block_turbo4_0 * __restrict__ dst, + const int64_t ne00, const int64_t ne01, const int64_t ne02, + const int64_t ne11, const int64_t ne12, + 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 int64_t s1, const int64_t s2, const int64_t s3, + const int innerq_slot) { + constexpr int GROUP_SIZE = QK_TURBO4; // 128 + const int j = threadIdx.x; + + const int64_t n_groups_per_row = ne00 / GROUP_SIZE; + const int64_t g = blockIdx.x; + const int64_t i_grp = g % n_groups_per_row; + int64_t tmp = g / n_groups_per_row; + const int64_t i01 = tmp % ne01; + tmp = tmp / ne01; + const int64_t i02 = tmp % ne02; + const int64_t i03 = tmp / ne02; + const int64_t i12 = i03 % ne12; + const int64_t i11 = i02 % ne11; + const int64_t i10 = i01; + + const int64_t dst_row = *(src1 + i10*s10 + i11*s11 + i12*s12); + const float * src_row = src0 + i01*s01 + i02*s02 + i03*s03; + block_turbo4_0 * dst_row_ptr = (block_turbo4_0 *)((char *)dst + dst_row*s1 + i02*s2 + i03*s3); + block_turbo4_0 * blk = dst_row_ptr + i_grp; + + __shared__ float x[GROUP_SIZE]; + x[j] = src_row[i_grp * GROUP_SIZE + j]; + __syncthreads(); + + // InnerQ calibrate/apply (per-width slot) — identical to turbo3. + const int innerq_hd = (innerq_slot >= 0) ? d_innerq_slot_hd[innerq_slot] : 0; + if (innerq_slot >= 0 && d_innerq_calibrating && innerq_hd > 0) { + const int ch = (int)((i_grp * GROUP_SIZE + j) % innerq_hd); + atomicAdd(&d_innerq_sq_accum[innerq_slot][ch], x[j] * x[j]); + if (j == 0) atomicAdd(&d_innerq_count, 1); + } + if (innerq_slot >= 0 && d_innerq_active && innerq_hd > 0) { + const int ch = (int)((i_grp * GROUP_SIZE + j) % innerq_hd); + x[j] *= d_innerq_scale[innerq_slot][ch]; + } + __syncthreads(); + + // L2 norm + constexpr int n_warps = GROUP_SIZE / WARP_SIZE; + __shared__ float warp_accum[n_warps]; + float v2 = x[j] * x[j]; + for (int offset = WARP_SIZE/2; offset > 0; offset >>= 1) v2 += __shfl_xor_sync(0xffffffff, v2, offset); + if (j % WARP_SIZE == 0) warp_accum[j / WARP_SIZE] = v2; + __syncthreads(); + __shared__ float s_norm_sq; + if (j == 0) { float t = 0.0f; for (int w = 0; w < n_warps; w++) t += warp_accum[w]; s_norm_sq = t; } + __syncthreads(); + const float grp_norm = sqrtf(s_norm_sq); + const float inv_norm = (grp_norm > 1e-10f) ? 1.0f / grp_norm : 0.0f; + + x[j] *= inv_norm; + __syncthreads(); + + // forward WHT + x[j] *= TURBO3_WHT_SIGNS1[j]; + __syncthreads(); +#define WHT_STAGE_T4(h) \ + if (j % (2*(h)) < (h)) { float a = x[j], b = x[j+(h)]; x[j] = a+b; x[j+(h)] = a-b; } \ + __syncthreads(); + WHT_STAGE_T4(1) WHT_STAGE_T4(2) WHT_STAGE_T4(4) + WHT_STAGE_T4(8) WHT_STAGE_T4(16) WHT_STAGE_T4(32) WHT_STAGE_T4(64) +#undef WHT_STAGE_T4 + const float inv_sqrt_group = 0.08838834764831845f; + x[j] = x[j] * inv_sqrt_group * TURBO3_WHT_SIGNS2[j]; + __syncthreads(); + + // quantize → 4-bit centroid index + const uint8_t idx = turbo4_nearest_centroid_4bit(x[j]); + + // pack: nibble, 2 elements per byte (low=even, high=odd) — no QJL signs + const int lane = j % WARP_SIZE; + const int elem_in_block = j; + const int qs_byte_idx = elem_in_block / 2; + const uint8_t my_nib = idx & 0xF; + uint8_t qs_byte = 0; +#pragma unroll + for (int k = 0; k < 2; k++) { + uint8_t contrib = __shfl_sync(0xffffffff, my_nib, (lane & ~1) + k); + qs_byte |= contrib << (k * 4); + } + if (lane % 2 == 0) blk->qs[qs_byte_idx] = qs_byte; + + // reconstruction norm + const float c = TURBO4_CENTROIDS_4BIT[idx]; + float rc = c * c; + for (int offset = WARP_SIZE/2; offset > 0; offset >>= 1) rc += __shfl_xor_sync(0xffffffff, rc, offset); + if (j % WARP_SIZE == 0) warp_accum[j / WARP_SIZE] = rc; + __syncthreads(); + __shared__ float s_recon_sq; + if (j == 0) { float t = 0.0f; for (int w = 0; w < n_warps; w++) t += warp_accum[w]; s_recon_sq = t; } + __syncthreads(); + const float recon_norm = sqrtf(s_recon_sq); + const float corrected_norm = (recon_norm > 1e-10f) ? grp_norm / recon_norm : grp_norm; + + if (elem_in_block == 0) { blk->norm = __float2half(corrected_norm); blk->rnorm = __float2half(0.0f); } +} + +template +static void set_rows_cuda_turbo4(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_TURBO4 == 0); + + cudaStream_t stream = ctx.stream(); + const int64_t s01 = nb01 / sizeof(float), s02 = nb02 / sizeof(float), s03 = nb03 / sizeof(float); + const int64_t s10 = nb10 / sizeof(idx_t), s11 = nb11 / sizeof(idx_t), s12 = nb12 / sizeof(idx_t); + const int64_t n_groups = ne00 / QK_TURBO4; + const int64_t ne_total = n_groups * ne01 * ne02 * ne03; + if (ne_total <= 0) return; + + turbo_innerq_check_finalize(QK_TURBO4); + const int innerq_slot = (innerq_mode != 0) ? innerq_slot_for((int) ne00) : -1; + + k_set_rows_turbo4<<<(int)ne_total, QK_TURBO4, 0, stream>>>( + src0_d, src1_d, (block_turbo4_0 *)dst->data, + ne00, ne01, ne02, ne11, ne12, + s01, s02, s03, s10, s11, s12, + nb1, nb2, nb3, innerq_slot); +} + +// ===================== turbo2 (2-bit, 4/byte) write kernel ===================== +// Clone of turbo3 minus QJL signs: 2-bit centroid + 4-per-byte pack (idx already 2-bit). +template +static __global__ void k_set_rows_turbo2( + const float * __restrict__ src0, + const idx_t * __restrict__ src1, + block_turbo2_0 * __restrict__ dst, + const int64_t ne00, const int64_t ne01, const int64_t ne02, + const int64_t ne11, const int64_t ne12, + 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 int64_t s1, const int64_t s2, const int64_t s3, + const int innerq_slot) { + constexpr int GROUP_SIZE = QK_TURBO2; // 128 + const int j = threadIdx.x; + + const int64_t n_groups_per_row = ne00 / GROUP_SIZE; + const int64_t g = blockIdx.x; + const int64_t i_grp = g % n_groups_per_row; + int64_t tmp = g / n_groups_per_row; + const int64_t i01 = tmp % ne01; + tmp = tmp / ne01; + const int64_t i02 = tmp % ne02; + const int64_t i03 = tmp / ne02; + const int64_t i12 = i03 % ne12; + const int64_t i11 = i02 % ne11; + const int64_t i10 = i01; + + const int64_t dst_row = *(src1 + i10*s10 + i11*s11 + i12*s12); + const float * src_row = src0 + i01*s01 + i02*s02 + i03*s03; + block_turbo2_0 * dst_row_ptr = (block_turbo2_0 *)((char *)dst + dst_row*s1 + i02*s2 + i03*s3); + block_turbo2_0 * blk = dst_row_ptr + i_grp; + + __shared__ float x[GROUP_SIZE]; + x[j] = src_row[i_grp * GROUP_SIZE + j]; + __syncthreads(); + + const int innerq_hd = (innerq_slot >= 0) ? d_innerq_slot_hd[innerq_slot] : 0; + if (innerq_slot >= 0 && d_innerq_calibrating && innerq_hd > 0) { + const int ch = (int)((i_grp * GROUP_SIZE + j) % innerq_hd); + atomicAdd(&d_innerq_sq_accum[innerq_slot][ch], x[j] * x[j]); + if (j == 0) atomicAdd(&d_innerq_count, 1); + } + if (innerq_slot >= 0 && d_innerq_active && innerq_hd > 0) { + const int ch = (int)((i_grp * GROUP_SIZE + j) % innerq_hd); + x[j] *= d_innerq_scale[innerq_slot][ch]; + } + __syncthreads(); + + constexpr int n_warps = GROUP_SIZE / WARP_SIZE; + __shared__ float warp_accum[n_warps]; + float v2 = x[j] * x[j]; + for (int offset = WARP_SIZE/2; offset > 0; offset >>= 1) v2 += __shfl_xor_sync(0xffffffff, v2, offset); + if (j % WARP_SIZE == 0) warp_accum[j / WARP_SIZE] = v2; + __syncthreads(); + __shared__ float s_norm_sq; + if (j == 0) { float t = 0.0f; for (int w = 0; w < n_warps; w++) t += warp_accum[w]; s_norm_sq = t; } + __syncthreads(); + const float grp_norm = sqrtf(s_norm_sq); + const float inv_norm = (grp_norm > 1e-10f) ? 1.0f / grp_norm : 0.0f; + + x[j] *= inv_norm; + __syncthreads(); + + x[j] *= TURBO3_WHT_SIGNS1[j]; + __syncthreads(); +#define WHT_STAGE_T2(h) \ + if (j % (2*(h)) < (h)) { float a = x[j], b = x[j+(h)]; x[j] = a+b; x[j+(h)] = a-b; } \ + __syncthreads(); + WHT_STAGE_T2(1) WHT_STAGE_T2(2) WHT_STAGE_T2(4) + WHT_STAGE_T2(8) WHT_STAGE_T2(16) WHT_STAGE_T2(32) WHT_STAGE_T2(64) +#undef WHT_STAGE_T2 + const float inv_sqrt_group = 0.08838834764831845f; + x[j] = x[j] * inv_sqrt_group * TURBO3_WHT_SIGNS2[j]; + __syncthreads(); + + const uint8_t idx = turbo2_nearest_centroid_2bit(x[j]); + + const int lane = j % WARP_SIZE; + const int elem_in_block = j; + const int qs_byte_idx = elem_in_block / 4; + const uint8_t my_2bit = idx & 0x3; + uint8_t qs_byte = 0; +#pragma unroll + for (int k = 0; k < 4; k++) { + uint8_t contrib = __shfl_sync(0xffffffff, my_2bit, (lane & ~3) + k); + qs_byte |= contrib << (k * 2); + } + if (lane % 4 == 0) blk->qs[qs_byte_idx] = qs_byte; + + const float c = TURBO2_CENTROIDS_2BIT[idx]; + float rc = c * c; + for (int offset = WARP_SIZE/2; offset > 0; offset >>= 1) rc += __shfl_xor_sync(0xffffffff, rc, offset); + if (j % WARP_SIZE == 0) warp_accum[j / WARP_SIZE] = rc; + __syncthreads(); + __shared__ float s_recon_sq; + if (j == 0) { float t = 0.0f; for (int w = 0; w < n_warps; w++) t += warp_accum[w]; s_recon_sq = t; } + __syncthreads(); + const float recon_norm = sqrtf(s_recon_sq); + const float corrected_norm = (recon_norm > 1e-10f) ? grp_norm / recon_norm : grp_norm; + + if (elem_in_block == 0) blk->norm = __float2half(corrected_norm); +} + +template +static void set_rows_cuda_turbo2(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 == 0); + + cudaStream_t stream = ctx.stream(); + const int64_t s01 = nb01 / sizeof(float), s02 = nb02 / sizeof(float), s03 = nb03 / sizeof(float); + const int64_t s10 = nb10 / sizeof(idx_t), s11 = nb11 / sizeof(idx_t), s12 = nb12 / sizeof(idx_t); + const int64_t n_groups = ne00 / QK_TURBO2; + const int64_t ne_total = n_groups * ne01 * ne02 * ne03; + if (ne_total <= 0) return; + + turbo_innerq_check_finalize(QK_TURBO2); + const int innerq_slot = (innerq_mode != 0) ? innerq_slot_for((int) ne00) : -1; + + k_set_rows_turbo2<<<(int)ne_total, QK_TURBO2, 0, stream>>>( + src0_d, src1_d, (block_turbo2_0 *)dst->data, + ne00, ne01, ne02, ne11, ne12, + s01, s02, s03, s10, s11, s12, + nb1, nb2, nb3, innerq_slot); +} + +// ===================== turbo8 (8-bit uniform int8) write kernel ===================== +// Clone of turbo3 preamble; quantize = uniform int8 (q=clamp(round(x*TURBO8_SCALE),-127,127)), +// pack = 1 byte/element (no warp packing), recon centroid = q/TURBO8_SCALE. +template +static __global__ void k_set_rows_turbo8( + const float * __restrict__ src0, + const idx_t * __restrict__ src1, + block_turbo8_0 * __restrict__ dst, + const int64_t ne00, const int64_t ne01, const int64_t ne02, + const int64_t ne11, const int64_t ne12, + 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 int64_t s1, const int64_t s2, const int64_t s3, + const int innerq_slot) { + constexpr int GROUP_SIZE = QK_TURBO8; // 128 + const int j = threadIdx.x; + + const int64_t n_groups_per_row = ne00 / GROUP_SIZE; + const int64_t g = blockIdx.x; + const int64_t i_grp = g % n_groups_per_row; + int64_t tmp = g / n_groups_per_row; + const int64_t i01 = tmp % ne01; + tmp = tmp / ne01; + const int64_t i02 = tmp % ne02; + const int64_t i03 = tmp / ne02; + const int64_t i12 = i03 % ne12; + const int64_t i11 = i02 % ne11; + const int64_t i10 = i01; + + const int64_t dst_row = *(src1 + i10*s10 + i11*s11 + i12*s12); + const float * src_row = src0 + i01*s01 + i02*s02 + i03*s03; + block_turbo8_0 * dst_row_ptr = (block_turbo8_0 *)((char *)dst + dst_row*s1 + i02*s2 + i03*s3); + block_turbo8_0 * blk = dst_row_ptr + i_grp; + + __shared__ float x[GROUP_SIZE]; + x[j] = src_row[i_grp * GROUP_SIZE + j]; + __syncthreads(); + + const int innerq_hd = (innerq_slot >= 0) ? d_innerq_slot_hd[innerq_slot] : 0; + if (innerq_slot >= 0 && d_innerq_calibrating && innerq_hd > 0) { + const int ch = (int)((i_grp * GROUP_SIZE + j) % innerq_hd); + atomicAdd(&d_innerq_sq_accum[innerq_slot][ch], x[j] * x[j]); + if (j == 0) atomicAdd(&d_innerq_count, 1); + } + if (innerq_slot >= 0 && d_innerq_active && innerq_hd > 0) { + const int ch = (int)((i_grp * GROUP_SIZE + j) % innerq_hd); + x[j] *= d_innerq_scale[innerq_slot][ch]; + } + __syncthreads(); + + constexpr int n_warps = GROUP_SIZE / WARP_SIZE; + __shared__ float warp_accum[n_warps]; + float v2 = x[j] * x[j]; + for (int offset = WARP_SIZE/2; offset > 0; offset >>= 1) v2 += __shfl_xor_sync(0xffffffff, v2, offset); + if (j % WARP_SIZE == 0) warp_accum[j / WARP_SIZE] = v2; + __syncthreads(); + __shared__ float s_norm_sq; + if (j == 0) { float t = 0.0f; for (int w = 0; w < n_warps; w++) t += warp_accum[w]; s_norm_sq = t; } + __syncthreads(); + const float grp_norm = sqrtf(s_norm_sq); + const float inv_norm = (grp_norm > 1e-10f) ? 1.0f / grp_norm : 0.0f; + + x[j] *= inv_norm; + __syncthreads(); + + x[j] *= TURBO3_WHT_SIGNS1[j]; + __syncthreads(); +#define WHT_STAGE_T8(h) \ + if (j % (2*(h)) < (h)) { float a = x[j], b = x[j+(h)]; x[j] = a+b; x[j+(h)] = a-b; } \ + __syncthreads(); + WHT_STAGE_T8(1) WHT_STAGE_T8(2) WHT_STAGE_T8(4) + WHT_STAGE_T8(8) WHT_STAGE_T8(16) WHT_STAGE_T8(32) WHT_STAGE_T8(64) +#undef WHT_STAGE_T8 + const float inv_sqrt_group = 0.08838834764831845f; + x[j] = x[j] * inv_sqrt_group * TURBO3_WHT_SIGNS2[j]; + __syncthreads(); + + // quantize → signed int8 (saturating), 1 byte/element + int q = __float2int_rn(x[j] * TURBO8_SCALE); + if (q > 127) q = 127; + if (q < -127) q = -127; + blk->qs[j] = (int8_t) q; + + const float c = (float) q / TURBO8_SCALE; + float rc = c * c; + for (int offset = WARP_SIZE/2; offset > 0; offset >>= 1) rc += __shfl_xor_sync(0xffffffff, rc, offset); + if (j % WARP_SIZE == 0) warp_accum[j / WARP_SIZE] = rc; + __syncthreads(); + __shared__ float s_recon_sq; + if (j == 0) { float t = 0.0f; for (int w = 0; w < n_warps; w++) t += warp_accum[w]; s_recon_sq = t; } + __syncthreads(); + const float recon_norm = sqrtf(s_recon_sq); + const float corrected_norm = (recon_norm > 1e-10f) ? grp_norm / recon_norm : grp_norm; + + if (j == 0) blk->norm = __float2half(corrected_norm); +} + +template +static void set_rows_cuda_turbo8(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_TURBO8 == 0); + + cudaStream_t stream = ctx.stream(); + const int64_t s01 = nb01 / sizeof(float), s02 = nb02 / sizeof(float), s03 = nb03 / sizeof(float); + const int64_t s10 = nb10 / sizeof(idx_t), s11 = nb11 / sizeof(idx_t), s12 = nb12 / sizeof(idx_t); + const int64_t n_groups = ne00 / QK_TURBO8; + const int64_t ne_total = n_groups * ne01 * ne02 * ne03; + if (ne_total <= 0) return; + + turbo_innerq_check_finalize(QK_TURBO8); + const int innerq_slot = (innerq_mode != 0) ? innerq_slot_for((int) ne00) : -1; + + k_set_rows_turbo8<<<(int)ne_total, QK_TURBO8, 0, stream>>>( + src0_d, src1_d, (block_turbo8_0 *)dst->data, + ne00, ne01, ne02, ne11, ne12, + s01, s02, s03, s10, s11, s12, + nb1, nb2, nb3, innerq_slot); +} + template static void set_rows_cuda(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { const src_t * src0_d = (const src_t *)src0->data; @@ -314,6 +1114,18 @@ static void set_rows_cuda(ggml_backend_cuda_context & ctx, const ggml_tensor * s nb1, nb2, nb3, stream ); + } else if (dst->type == GGML_TYPE_TURBO3_0) { + // opencoti-hook: turboquant — see docs/features/poly_kv.md (M6-S2 Gate-2 write path) + set_rows_cuda_turbo3(ctx, src0, src1, dst); + } else if (dst->type == GGML_TYPE_TURBO4_0) { + // opencoti-hook: turboquant tier family — see docs/features/poly_kv.md (M6-S2) + set_rows_cuda_turbo4(ctx, src0, src1, dst); + } else if (dst->type == GGML_TYPE_TURBO2_0) { + // opencoti-hook: turboquant tier family — see docs/features/poly_kv.md (M6-S2) + set_rows_cuda_turbo2(ctx, src0, src1, dst); + } 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 { GGML_ABORT("unsupported type %s", ggml_type_name(dst->type)); } diff --git a/llama.cpp/ggml/src/ggml-cuda/template-instances/fattn-vec-instance-turbo2_0-turbo2_0.cu b/llama.cpp/ggml/src/ggml-cuda/template-instances/fattn-vec-instance-turbo2_0-turbo2_0.cu new file mode 100644 --- /dev/null +++ b/llama.cpp/ggml/src/ggml-cuda/template-instances/fattn-vec-instance-turbo2_0-turbo2_0.cu @@ -0,0 +1,8 @@ +// This file has been autogenerated by generate_cu_files.py, do not edit manually. +// opencoti-hook: turboquant perf-lift (Level B) — F5 M6-S2 #392 + +#include "../fattn-vec.cuh" + +DECL_FATTN_VEC_CASE( 64, GGML_TYPE_TURBO2_0, GGML_TYPE_TURBO2_0); +DECL_FATTN_VEC_CASE(128, GGML_TYPE_TURBO2_0, GGML_TYPE_TURBO2_0); +DECL_FATTN_VEC_CASE(256, GGML_TYPE_TURBO2_0, GGML_TYPE_TURBO2_0); diff --git a/llama.cpp/ggml/src/ggml-cuda/template-instances/fattn-vec-instance-turbo3_0-turbo3_0.cu b/llama.cpp/ggml/src/ggml-cuda/template-instances/fattn-vec-instance-turbo3_0-turbo3_0.cu new file mode 100644 --- /dev/null +++ b/llama.cpp/ggml/src/ggml-cuda/template-instances/fattn-vec-instance-turbo3_0-turbo3_0.cu @@ -0,0 +1,8 @@ +// This file has been autogenerated by generate_cu_files.py, do not edit manually. +// opencoti-hook: turboquant perf-lift (Level B) — F5 M6-S2 #392 + +#include "../fattn-vec.cuh" + +DECL_FATTN_VEC_CASE( 64, GGML_TYPE_TURBO3_0, GGML_TYPE_TURBO3_0); +DECL_FATTN_VEC_CASE(128, GGML_TYPE_TURBO3_0, GGML_TYPE_TURBO3_0); +DECL_FATTN_VEC_CASE(256, GGML_TYPE_TURBO3_0, GGML_TYPE_TURBO3_0); diff --git a/llama.cpp/ggml/src/ggml-cuda/template-instances/fattn-vec-instance-turbo4_0-turbo4_0.cu b/llama.cpp/ggml/src/ggml-cuda/template-instances/fattn-vec-instance-turbo4_0-turbo4_0.cu new file mode 100644 --- /dev/null +++ b/llama.cpp/ggml/src/ggml-cuda/template-instances/fattn-vec-instance-turbo4_0-turbo4_0.cu @@ -0,0 +1,8 @@ +// This file has been autogenerated by generate_cu_files.py, do not edit manually. +// opencoti-hook: turboquant perf-lift (Level B) — F5 M6-S2 #392 + +#include "../fattn-vec.cuh" + +DECL_FATTN_VEC_CASE( 64, GGML_TYPE_TURBO4_0, GGML_TYPE_TURBO4_0); +DECL_FATTN_VEC_CASE(128, GGML_TYPE_TURBO4_0, GGML_TYPE_TURBO4_0); +DECL_FATTN_VEC_CASE(256, GGML_TYPE_TURBO4_0, GGML_TYPE_TURBO4_0); diff --git a/llama.cpp/ggml/src/ggml-cuda/template-instances/generate_cu_files.py b/llama.cpp/ggml/src/ggml-cuda/template-instances/generate_cu_files.py --- a/llama.cpp/ggml/src/ggml-cuda/template-instances/generate_cu_files.py +++ b/llama.cpp/ggml/src/ggml-cuda/template-instances/generate_cu_files.py @@ -74,6 +74,12 @@ for type_k in TYPES_KV: with open(f"fattn-vec-instance-{get_short_name(type_k)}-{get_short_name(type_v)}.cu", "w") as f: f.write(SOURCE_FATTN_VEC.format(type_k=type_k, type_v=type_v)) +# opencoti-hook: turboquant perf-lift (Level B) — F5 M6-S2 #392. Fused turbo K/V vec instances +# (K==V same tier only; the in-register readers live in fattn-turbo.cuh). +for type_t in ["GGML_TYPE_TURBO2_0", "GGML_TYPE_TURBO3_0", "GGML_TYPE_TURBO4_0"]: + with open(f"fattn-vec-instance-{get_short_name(type_t)}-{get_short_name(type_t)}.cu", "w") as f: + f.write(SOURCE_FATTN_VEC.format(type_k=type_t, type_v=type_t)) + for ncols in [8, 16, 32, 64]: for ncols2 in [1, 2, 4, 8, 16, 32]: if ncols2 > ncols: diff --git a/llama.cpp/ggml/src/ggml-cuda/turbo-cpy.cu b/llama.cpp/ggml/src/ggml-cuda/turbo-cpy.cu new file mode 100644 --- /dev/null +++ b/llama.cpp/ggml/src/ggml-cuda/turbo-cpy.cu @@ -0,0 +1,133 @@ +// opencoti-hook: turboquant perf-lift (M6-S2 Level A) — see docs/features/poly_kv.md +// +// GPU-resident dequant-on-lift kernels for TurboQuant KV (turbo2/3/4/8 → f16). One CUDA threadblock +// (128 threads) per 128-element WHT group: load codebook/int8 → centroid×norm into shared mem → +// inverse WHT computed ONCE cooperatively (not per output pair) → ×scale_inv[ch] (InnerQ, true +// space) → strided f16 write. Numerically mirrors the CPU dequant dequantize_row_turboN_0 +// (ggml-turbo-quant.c): same inverse-WHT order (signs2 → butterfly → 1/√128·signs1), same per-width +// channel fold (global_index % head_dim), same InnerQ un-equalize. Accept = logit-equivalence. + +#include "common.cuh" +#include "turbo-dequant.cuh" // TURBO{3}_WHT_SIGNS1/2 + TURBO{2,3,4}_CENTROIDS tables (this TU's copy) +#include "turbo-cpy.cuh" + +// Cooperative inverse WHT over shared buf[128] (already holding centroid×norm in rotated space), +// then InnerQ ×scale_inv and a strided f16 write of element e0+t. All 128 threads of the block +// participate; callers guarantee e0+t < ne (no divergent early-return before the syncs). +static __device__ __forceinline__ void turbo_wht_inv_write( + float * buf, int t, int64_t e0, char * cdst, + int64_t ne10, int64_t ne11, int64_t ne12, + int64_t nb10, int64_t nb11, int64_t nb12, int64_t nb13, + const float * scale_inv, int head_dim) { + __syncthreads(); + buf[t] *= TURBO3_WHT_SIGNS2[t]; + for (int h = 1; h < 128; h <<= 1) { + __syncthreads(); + if ((t & h) == 0) { + const float a = buf[t], b = buf[t + h]; + buf[t] = a + b; + buf[t + h] = a - b; + } + } + __syncthreads(); + float val = buf[t] * 0.08838834764831845f * TURBO3_WHT_SIGNS1[t]; + if (scale_inv) { + val *= scale_inv[(int) ((e0 + t) % head_dim)]; + } + const int64_t e = e0 + t; + const int64_t i13 = e / (ne10 * ne11 * ne12); + const int64_t i12 = (e - i13*ne10*ne11*ne12) / (ne10*ne11); + const int64_t i11 = (e - i13*ne10*ne11*ne12 - i12*ne10*ne11) / ne10; + const int64_t i10 = e - i13*ne10*ne11*ne12 - i12*ne10*ne11 - i11*ne10; + const int64_t dst_off = i10*nb10 + i11*nb11 + i12*nb12 + i13*nb13; + *(half *) (cdst + dst_off) = __float2half(val); +} + +// Compute the (shared) turbo block pointer for group e0 from the strided src layout. +#define TURBO_SRC_BLOCK(block_t) \ + const int64_t i03 = e0 / (ne00 * ne01 * ne02); \ + const int64_t i02 = (e0 - i03*ne00*ne01*ne02) / (ne00*ne01); \ + const int64_t i01 = (e0 - i03*ne00*ne01*ne02 - i02*ne01*ne00) / ne00; \ + const int64_t i00 = e0 - i03*ne00*ne01*ne02 - i02*ne01*ne00 - i01*ne00; \ + const int64_t x_off = (i00 / 128)*nb00 + i01*nb01 + i02*nb02 + i03*nb03; \ + const block_t * blk = (const block_t *) (cx + x_off) + +#define TURBO_CPY_PARAMS \ + const char * cx, char * cdst, int64_t ne, \ + int64_t ne00, int64_t ne01, int64_t ne02, \ + int64_t nb00, int64_t nb01, int64_t nb02, int64_t nb03, \ + int64_t ne10, int64_t ne11, int64_t ne12, \ + int64_t nb10, int64_t nb11, int64_t nb12, int64_t nb13, \ + const float * scale_inv, int head_dim + +static __global__ void cpy_turbo3_f16(TURBO_CPY_PARAMS) { + const int t = threadIdx.x; + const int64_t e0 = (int64_t) blockIdx.x * 128; + if (e0 >= ne) return; + TURBO_SRC_BLOCK(block_turbo3_0); + const float norm = __half2float(blk->norm); + __shared__ float buf[128]; + const uint8_t low2 = (blk->qs[t / 4] >> ((t % 4) * 2)) & 0x3; + const uint8_t hi1 = (blk->signs[t / 8] >> (t % 8)) & 0x1; + buf[t] = TURBO3_CENTROIDS_3BIT[(uint8_t)(low2 | (hi1 << 2))] * norm; + turbo_wht_inv_write(buf, t, e0, cdst, ne10, ne11, ne12, nb10, nb11, nb12, nb13, scale_inv, head_dim); +} + +static __global__ void cpy_turbo4_f16(TURBO_CPY_PARAMS) { + const int t = threadIdx.x; + const int64_t e0 = (int64_t) blockIdx.x * 128; + if (e0 >= ne) return; + TURBO_SRC_BLOCK(block_turbo4_0); + const float norm = __half2float(blk->norm); + __shared__ float buf[128]; + const uint8_t idx = (blk->qs[t / 2] >> ((t % 2) * 4)) & 0xF; + buf[t] = TURBO4_CENTROIDS_4BIT[idx] * norm; + turbo_wht_inv_write(buf, t, e0, cdst, ne10, ne11, ne12, nb10, nb11, nb12, nb13, scale_inv, head_dim); +} + +static __global__ void cpy_turbo2_f16(TURBO_CPY_PARAMS) { + const int t = threadIdx.x; + const int64_t e0 = (int64_t) blockIdx.x * 128; + if (e0 >= ne) return; + TURBO_SRC_BLOCK(block_turbo2_0); + const float norm = __half2float(blk->norm); + __shared__ float buf[128]; + const uint8_t idx = (blk->qs[t / 4] >> ((t % 4) * 2)) & 0x3; + buf[t] = TURBO2_CENTROIDS_2BIT[idx] * norm; + turbo_wht_inv_write(buf, t, e0, cdst, ne10, ne11, ne12, nb10, nb11, nb12, nb13, scale_inv, head_dim); +} + +static __global__ void cpy_turbo8_f16(TURBO_CPY_PARAMS) { + const int t = threadIdx.x; + const int64_t e0 = (int64_t) blockIdx.x * 128; + if (e0 >= ne) return; + TURBO_SRC_BLOCK(block_turbo8_0); + const float norm = __half2float(blk->norm); + __shared__ float buf[128]; + buf[t] = (float) blk->qs[t] * (1.0f / TURBO8_SCALE) * norm; + turbo_wht_inv_write(buf, t, e0, cdst, ne10, ne11, ne12, nb10, nb11, nb12, nb13, scale_inv, head_dim); +} + +void ggml_cpy_turbo_f16_cuda( + enum ggml_type src_type, const char * cx, char * cdst, int64_t ne, + int64_t ne00, int64_t ne01, int64_t ne02, + int64_t nb00, int64_t nb01, int64_t nb02, int64_t nb03, + int64_t ne10, int64_t ne11, int64_t ne12, + int64_t nb10, int64_t nb11, int64_t nb12, int64_t nb13, + const float * d_scale_inv, int head_dim, cudaStream_t stream) { + GGML_ASSERT(ne % 128 == 0); + const int64_t num_groups = ne / 128; + GGML_ASSERT(num_groups < UINT_MAX); + const dim3 grid((unsigned) num_groups); + #define LAUNCH(K) K<<>>( \ + cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, \ + ne10, ne11, ne12, nb10, nb11, nb12, nb13, d_scale_inv, head_dim) + switch (src_type) { + case GGML_TYPE_TURBO3_0: LAUNCH(cpy_turbo3_f16); break; + case GGML_TYPE_TURBO4_0: LAUNCH(cpy_turbo4_f16); break; + case GGML_TYPE_TURBO2_0: LAUNCH(cpy_turbo2_f16); break; + case GGML_TYPE_TURBO8_0: LAUNCH(cpy_turbo8_f16); break; + default: GGML_ABORT("ggml_cpy_turbo_f16_cuda: non-turbo src_type %d\n", (int) src_type); + } + #undef LAUNCH +} diff --git a/llama.cpp/ggml/src/ggml-cuda/turbo-cpy.cuh b/llama.cpp/ggml/src/ggml-cuda/turbo-cpy.cuh new file mode 100644 --- /dev/null +++ b/llama.cpp/ggml/src/ggml-cuda/turbo-cpy.cuh @@ -0,0 +1,25 @@ +// opencoti-hook: turboquant perf-lift (M6-S2 Level A) — see docs/features/poly_kv.md +// +// GPU-resident dequant-on-lift for TurboQuant KV. Replaces the host round-trip (the scheduler used +// to spill ggml_cast(turboN→f16) to the CPU backend because CUDA's CPY supports_op rejected turbo) +// with a single block-cooperative kernel: one CUDA threadblock per 128-element WHT group computes +// the inverse WHT ONCE (vs the per-output-pair recompute in turbo-dequant.cuh's convert path) and +// applies InnerQ ×scale_inv in true-channel space. Output is f16 in VRAM; stock flash-attention +// then runs unchanged. Numerically reproduces the CPU dequant (dequantize_row_turboN_0); gated by +// logit-equivalence, never byte-equality. +#pragma once + +#include "ggml.h" +#include + +// Dispatch a turboN→f16 dequant-cpy on the GPU. src_type ∈ {TURBO2_0,TURBO3_0,TURBO4_0,TURBO8_0}. +// d_scale_inv: device pointer to the active InnerQ slot's scale_inv[head_dim] (from +// turbo_innerq_active_device_scale), or nullptr ⇒ no un-equalize (identity). ne must be a multiple +// of 128 (the WHT group / turbo block size). +void ggml_cpy_turbo_f16_cuda( + enum ggml_type src_type, const char * cx, char * cdst, int64_t ne, + int64_t ne00, int64_t ne01, int64_t ne02, + int64_t nb00, int64_t nb01, int64_t nb02, int64_t nb03, + int64_t ne10, int64_t ne11, int64_t ne12, + int64_t nb10, int64_t nb11, int64_t nb12, int64_t nb13, + const float * d_scale_inv, int head_dim, cudaStream_t stream); diff --git a/llama.cpp/ggml/src/ggml-cuda/turbo-dequant.cuh b/llama.cpp/ggml/src/ggml-cuda/turbo-dequant.cuh new file mode 100644 --- /dev/null +++ b/llama.cpp/ggml/src/ggml-cuda/turbo-dequant.cuh @@ -0,0 +1,213 @@ +// opencoti-hook: turboquant — see docs/features/poly_kv.md (M6-S2 Gate-2) +// +// Minimal, InnerQ-FREE CUDA dequant for TurboQuant turbo3 KV, used by the +// Gate-2 "dequant-then-standard-FA" fallback. Folds the INVERSE WHT into the +// block dequant so the emitted f16 is in TRUE (un-rotated) space — the +// standard f16 flash-attention then computes q·k directly with NO q-rotation +// hook. (The fork's hot fused kernels instead store rotated K and rotate q; +// that path + InnerQ Q/V-equalization is Gate-3.) +// +// turbo3 block = 128 values = exactly one WHT group, so the inverse WHT is a +// self-contained per-block 128-pt butterfly (cold path: recomputed per element +// pair, like the fork's dequantize_tq3_1s). Constants copied verbatim from the +// AtomicBot fork ggml-cuda/turbo-quant.cuh; inverse-WHT order mirrors the +// cosine-0.983-validated turbo_cpu_fwht_inverse (signs2 → butterfly → 1/√128·signs1). +#pragma once + +#include "common.cuh" + +#define QK_TURBO3_INV 128 + +static __constant__ float TURBO3_CENTROIDS_3BIT[8] = { + -0.190685f, -0.117832f, -0.065717f, -0.021460f, + 0.021460f, 0.065717f, 0.117832f, 0.190685f +}; + +static __constant__ float TURBO3_WHT_SIGNS1[128] = { + -1.0f, 1.0f, 1.0f, -1.0f, -1.0f, 1.0f, -1.0f, 1.0f, -1.0f, -1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, + 1.0f, -1.0f, 1.0f, -1.0f, 1.0f, -1.0f, -1.0f, 1.0f, 1.0f, 1.0f, -1.0f, 1.0f, 1.0f, -1.0f, -1.0f, -1.0f, + -1.0f, 1.0f, 1.0f, -1.0f, 1.0f, 1.0f, -1.0f, 1.0f, -1.0f, 1.0f, 1.0f, -1.0f, -1.0f, 1.0f, -1.0f, 1.0f, + 1.0f, 1.0f, 1.0f, -1.0f, -1.0f, -1.0f, -1.0f, -1.0f, 1.0f, -1.0f, 1.0f, 1.0f, 1.0f, 1.0f, -1.0f, 1.0f, + -1.0f, -1.0f, 1.0f, -1.0f, -1.0f, -1.0f, 1.0f, -1.0f, -1.0f, -1.0f, 1.0f, -1.0f, -1.0f, -1.0f, 1.0f, 1.0f, + 1.0f, -1.0f, -1.0f, 1.0f, 1.0f, 1.0f, -1.0f, -1.0f, 1.0f, 1.0f, -1.0f, 1.0f, 1.0f, -1.0f, 1.0f, -1.0f, + -1.0f, 1.0f, 1.0f, -1.0f, 1.0f, -1.0f, 1.0f, -1.0f, 1.0f, 1.0f, 1.0f, 1.0f, -1.0f, 1.0f, -1.0f, 1.0f, + 1.0f, -1.0f, 1.0f, 1.0f, -1.0f, -1.0f, -1.0f, -1.0f, -1.0f, 1.0f, 1.0f, -1.0f, 1.0f, 1.0f, -1.0f, 1.0f +}; + +static __constant__ float TURBO3_WHT_SIGNS2[128] = { + 1.0f, 1.0f, 1.0f, 1.0f, -1.0f, 1.0f, 1.0f, -1.0f, 1.0f, -1.0f, -1.0f, -1.0f, 1.0f, -1.0f, -1.0f, -1.0f, + 1.0f, 1.0f, -1.0f, -1.0f, 1.0f, -1.0f, 1.0f, -1.0f, 1.0f, -1.0f, -1.0f, 1.0f, -1.0f, 1.0f, 1.0f, 1.0f, + 1.0f, 1.0f, -1.0f, -1.0f, -1.0f, 1.0f, -1.0f, -1.0f, -1.0f, -1.0f, -1.0f, -1.0f, 1.0f, 1.0f, 1.0f, -1.0f, + 1.0f, -1.0f, 1.0f, 1.0f, 1.0f, -1.0f, -1.0f, 1.0f, -1.0f, -1.0f, -1.0f, -1.0f, -1.0f, -1.0f, 1.0f, 1.0f, + 1.0f, -1.0f, 1.0f, -1.0f, -1.0f, -1.0f, -1.0f, 1.0f, -1.0f, 1.0f, -1.0f, 1.0f, -1.0f, -1.0f, 1.0f, 1.0f, + -1.0f, 1.0f, -1.0f, 1.0f, 1.0f, -1.0f, 1.0f, -1.0f, -1.0f, -1.0f, -1.0f, 1.0f, -1.0f, -1.0f, 1.0f, -1.0f, + 1.0f, -1.0f, 1.0f, 1.0f, 1.0f, -1.0f, -1.0f, 1.0f, -1.0f, 1.0f, -1.0f, 1.0f, 1.0f, -1.0f, -1.0f, 1.0f, + -1.0f, 1.0f, -1.0f, 1.0f, 1.0f, -1.0f, 1.0f, -1.0f, 1.0f, -1.0f, -1.0f, -1.0f, -1.0f, -1.0f, 1.0f, -1.0f +}; + +// ---- Forward-quantize side (used by set-rows.cu k_set_rows_turbo3, M6-S2 Gate-2 write path) ---- +static __constant__ float TURBO3_MID_3BIT[7] = { + -0.154259f, -0.091775f, -0.043589f, 0.0f, + 0.043589f, 0.091775f, 0.154259f +}; + +static __device__ __forceinline__ uint8_t turbo3_nearest_centroid_3bit(float val) { + if (val < TURBO3_MID_3BIT[0]) return 0; + else if (val < TURBO3_MID_3BIT[1]) return 1; + else if (val < TURBO3_MID_3BIT[2]) return 2; + else if (val < TURBO3_MID_3BIT[3]) return 3; + else if (val < TURBO3_MID_3BIT[4]) return 4; + else if (val < TURBO3_MID_3BIT[5]) return 5; + else if (val < TURBO3_MID_3BIT[6]) return 6; + else return 7; +} + +// Dequant one turbo3 block (128 vals) into TRUE space: centroid lookup → inverse WHT. +// Cold path — the full 128-pt butterfly is recomputed per (iqs) pair, matching the +// fork's dequantize_tq3_1s cold-dequant pattern. Correctness over speed (Gate-3 fuses). +static __device__ __forceinline__ void dequantize_turbo3_0_inv( + const void * vx, const int64_t ib, const int iqs, float2 & v) { + const block_turbo3_0 * x = (const block_turbo3_0 *) vx; + const float norm = __half2float(x[ib].norm); + + float buf[128]; + // centroid lookup (rotated space) + for (int j = 0; j < 128; j++) { + const uint8_t low2 = (x[ib].qs[j / 4] >> ((j % 4) * 2)) & 0x3; + const uint8_t hi1 = (x[ib].signs[j / 8] >> (j % 8)) & 0x1; + const uint8_t idx = (uint8_t) (low2 | (hi1 << 2)); + buf[j] = TURBO3_CENTROIDS_3BIT[idx] * norm; + } + // inverse WHT: signs2 → butterfly → (1/sqrt128)·signs1 (validated order) + for (int j = 0; j < 128; j++) buf[j] *= TURBO3_WHT_SIGNS2[j]; + for (int h = 1; h < 128; h <<= 1) { + for (int i = 0; i < 128; i += h << 1) { + for (int j = i; j < i + h; j++) { + const float a = buf[j], b = buf[j + h]; + buf[j] = a + b; buf[j + h] = a - b; + } + } + } + const float inv_sqrt128 = 0.08838834764831845f; + for (int j = 0; j < 128; j++) buf[j] *= inv_sqrt128 * TURBO3_WHT_SIGNS1[j]; + + v.x = buf[iqs + 0]; + v.y = buf[iqs + 1]; +} + +// ===================== turbo4 (4-bit nibble PolarQuant) ===================== +// Same WHT group (128) ⇒ reuse TURBO3_WHT_SIGNS1/2 (width-independent). 16 MSE-optimal +// centroids for N(0,1/128), nibble-packed, no QJL signs. Centroids + midpoints mirror the +// CPU nearest_centroid_4bit / CENTROIDS_4BIT so GPU-write ↔ CPU-dequant indices agree. +static __constant__ float TURBO4_CENTROIDS_4BIT[16] = { + -0.173926f, -0.117195f, -0.089527f, -0.068756f, + -0.051262f, -0.035597f, -0.020989f, -0.006938f, + 0.006938f, 0.020989f, 0.035597f, 0.051262f, + 0.068756f, 0.089527f, 0.117195f, 0.173926f +}; +static __constant__ float TURBO4_MID_4BIT[15] = { + -0.145560f, -0.103361f, -0.079142f, -0.060009f, + -0.043430f, -0.028293f, -0.013963f, 0.0f, + 0.013963f, 0.028293f, 0.043430f, 0.060009f, + 0.079142f, 0.103361f, 0.145560f +}; +static __device__ __forceinline__ uint8_t turbo4_nearest_centroid_4bit(float val) { + if (val < TURBO4_MID_4BIT[0]) return 0; + else if (val < TURBO4_MID_4BIT[1]) return 1; + else if (val < TURBO4_MID_4BIT[2]) return 2; + else if (val < TURBO4_MID_4BIT[3]) return 3; + else if (val < TURBO4_MID_4BIT[4]) return 4; + else if (val < TURBO4_MID_4BIT[5]) return 5; + else if (val < TURBO4_MID_4BIT[6]) return 6; + else if (val < TURBO4_MID_4BIT[7]) return 7; + else if (val < TURBO4_MID_4BIT[8]) return 8; + else if (val < TURBO4_MID_4BIT[9]) return 9; + else if (val < TURBO4_MID_4BIT[10]) return 10; + else if (val < TURBO4_MID_4BIT[11]) return 11; + else if (val < TURBO4_MID_4BIT[12]) return 12; + else if (val < TURBO4_MID_4BIT[13]) return 13; + else if (val < TURBO4_MID_4BIT[14]) return 14; + else return 15; +} +static __device__ __forceinline__ void dequantize_turbo4_0_inv( + const void * vx, const int64_t ib, const int iqs, float2 & v) { + const block_turbo4_0 * x = (const block_turbo4_0 *) vx; + const float norm = __half2float(x[ib].norm); + float buf[128]; + for (int j = 0; j < 128; j++) { + const uint8_t idx = (x[ib].qs[j / 2] >> ((j % 2) * 4)) & 0xF; + buf[j] = TURBO4_CENTROIDS_4BIT[idx] * norm; + } + for (int j = 0; j < 128; j++) buf[j] *= TURBO3_WHT_SIGNS2[j]; + for (int h = 1; h < 128; h <<= 1) { + for (int i = 0; i < 128; i += h << 1) { + for (int j = i; j < i + h; j++) { + const float a = buf[j], b = buf[j + h]; + buf[j] = a + b; buf[j + h] = a - b; + } + } + } + const float inv_sqrt128 = 0.08838834764831845f; + for (int j = 0; j < 128; j++) buf[j] *= inv_sqrt128 * TURBO3_WHT_SIGNS1[j]; + v.x = buf[iqs + 0]; + v.y = buf[iqs + 1]; +} + +// ===================== turbo8 (8-bit uniform int8) ===================== +// q/TURBO8_SCALE → inverse WHT (no codebook). Reuses TURBO3_WHT_SIGNS1/2 (width-independent). +static __device__ __forceinline__ void dequantize_turbo8_0_inv( + const void * vx, const int64_t ib, const int iqs, float2 & v) { + const block_turbo8_0 * x = (const block_turbo8_0 *) vx; + const float norm = __half2float(x[ib].norm); + const float inv_scale = 1.0f / TURBO8_SCALE; + float buf[128]; + for (int j = 0; j < 128; j++) buf[j] = (float) x[ib].qs[j] * inv_scale * norm; + for (int j = 0; j < 128; j++) buf[j] *= TURBO3_WHT_SIGNS2[j]; + for (int h = 1; h < 128; h <<= 1) { + for (int i = 0; i < 128; i += h << 1) { + for (int j = i; j < i + h; j++) { + const float a = buf[j], b = buf[j + h]; + buf[j] = a + b; buf[j + h] = a - b; + } + } + } + const float inv_sqrt128 = 0.08838834764831845f; + for (int j = 0; j < 128; j++) buf[j] *= inv_sqrt128 * TURBO3_WHT_SIGNS1[j]; + v.x = buf[iqs + 0]; + v.y = buf[iqs + 1]; +} + +// ===================== turbo2 (2-bit PolarQuant, no QJL) ===================== +static __constant__ float TURBO2_CENTROIDS_2BIT[4] = { + -0.133462f, -0.039994f, 0.039994f, 0.133462f +}; +static __constant__ float TURBO2_MID_2BIT[3] = { -0.086728f, 0.0f, 0.086728f }; +static __device__ __forceinline__ uint8_t turbo2_nearest_centroid_2bit(float val) { + if (val < TURBO2_MID_2BIT[0]) return 0; + else if (val < TURBO2_MID_2BIT[1]) return 1; + else if (val < TURBO2_MID_2BIT[2]) return 2; + else return 3; +} +static __device__ __forceinline__ void dequantize_turbo2_0_inv( + const void * vx, const int64_t ib, const int iqs, float2 & v) { + const block_turbo2_0 * x = (const block_turbo2_0 *) vx; + const float norm = __half2float(x[ib].norm); + float buf[128]; + for (int j = 0; j < 128; j++) { + const uint8_t idx = (x[ib].qs[j / 4] >> ((j % 4) * 2)) & 0x3; + buf[j] = TURBO2_CENTROIDS_2BIT[idx] * norm; + } + for (int j = 0; j < 128; j++) buf[j] *= TURBO3_WHT_SIGNS2[j]; + for (int h = 1; h < 128; h <<= 1) { + for (int i = 0; i < 128; i += h << 1) { + for (int j = i; j < i + h; j++) { + const float a = buf[j], b = buf[j + h]; + buf[j] = a + b; buf[j + h] = a - b; + } + } + } + const float inv_sqrt128 = 0.08838834764831845f; + for (int j = 0; j < 128; j++) buf[j] *= inv_sqrt128 * TURBO3_WHT_SIGNS1[j]; + v.x = buf[iqs + 0]; + v.y = buf[iqs + 1]; +} diff --git a/llama.cpp/ggml/src/ggml-cuda/turbo-innerq-dev.cuh b/llama.cpp/ggml/src/ggml-cuda/turbo-innerq-dev.cuh new file mode 100644 --- /dev/null +++ b/llama.cpp/ggml/src/ggml-cuda/turbo-innerq-dev.cuh @@ -0,0 +1,21 @@ +// opencoti-hook: turboquant perf-lift (M6-S2 Level A) — see docs/features/poly_kv.md +// +// Cross-TU bridge for the GPU-resident dequant-on-lift. The InnerQ per-width scale_inv lives as +// TU-local `static __device__` symbols in set-rows.cu (d_innerq_scale_inv) and is therefore not +// reachable from another translation unit. This accessor — implemented in set-rows.cu, declared +// here, called by turbo-cpy.cu — exposes the device pointer to the *active* slot's scale_inv row +// plus its head_dim, selected by the cast's row width (ne00 = n_embd_gqa). It mirrors the CPU +// dequant's slot selection (turbo_innerq_cpu_slot_for in ggml-turbo-quant.c) so GPU lift and CPU +// lift apply identical InnerQ un-equalization. +// +// Returns false when InnerQ is inactive (default-off, calibrating, or no slot matches this width) — +// the caller then performs the inverse-WHT dequant with NO ×scale_inv (identity), exactly as the +// CPU path does when its slot lookup misses. +#pragma once + +#include + +// d_scale_inv_out: device pointer to scale_inv[head_dim] for the matched width (valid while the +// process lives; the symbol is static-storage device memory). head_dim_out: that slot's head_dim +// (the channel fold modulus). Both written only when the function returns true. +bool turbo_innerq_active_device_scale(int64_t ne00, const float ** d_scale_inv_out, int * head_dim_out); diff --git a/llama.cpp/ggml/src/ggml-cuda/turbo-innerq.cuh b/llama.cpp/ggml/src/ggml-cuda/turbo-innerq.cuh new file mode 100644 --- /dev/null +++ b/llama.cpp/ggml/src/ggml-cuda/turbo-innerq.cuh @@ -0,0 +1,114 @@ +// opencoti-hook: turboquant InnerQ (M6-S2) — see docs/features/poly_kv.md +// +// TurboQuant InnerQ per-channel equalization — calibration artifact I/O. +// +// InnerQ equalizes per-channel K variance BEFORE the WHT rotation so the shared +// 3-bit centroid codebook fits each channel: forward does x[j] *= scale[j], the +// dequant undoes it with x[j] *= scale_inv[j] in true-channel space. This is the +// missing correctness ingredient that makes turbo3 usable as a K-quant (turbo3-K +// at strength-0 is only ~0.983 cosine — too lossy for softmax attention scores). +// +// opencoti uses a CALIBRATE-then-ACTIVATE-FROM-TOKEN-0 flow so the whole KV cache +// is uniformly equalized (no mixed early/late tokens). Because the turbo3 write is +// GPU-only (this DSO) but our active dequant is the CPU one (host binary), the +// scale_inv[128] vector crosses the host/DSO boundary as a FILE ARTIFACT both +// sides load at boot: +// TURBO_INNERQ_CALIBRATE=N + TURBO_INNERQ_OUT=path : gather K² over N blocks, +// finalize, write `path`. (Cache discarded.) +// TURBO_INNERQ_SCALE=path : load scale_inv at init, active from token 0. +// Header-only (static-inline) so it adds no .cu to the build; the device state and +// init/finalize live in set-rows.cu (the only TU needing device access). +#pragma once + +#include +#include + +#define INNERQ_MAX_CHANNELS 512 // max head_dim (Gemma-4 global = 512); Qwen = 128, Gemma SWA = 256 +#define INNERQ_MAX_SLOTS 4 // distinct KV widths per model (Gemma-4 = 2: SWA+global) +#define TURBO_INNERQ_MAGIC 0x31514954u // "TIQ1" little-endian (v1: single uniform head_dim) +#define TURBO_INNERQ_MAGIC2 0x32514954u // "TIQ2" little-endian (v2: per-width slots) + +// File layout: u32 magic, i32 count, float[count] scale_inv (host-native LE). +static inline bool turbo_innerq_dump_file(const char * path, const float * scale_inv, int n) { + if (!path || n <= 0 || n > INNERQ_MAX_CHANNELS) return false; + FILE * f = fopen(path, "wb"); + if (!f) return false; + const uint32_t magic = TURBO_INNERQ_MAGIC; + const int32_t cnt = n; + bool ok = fwrite(&magic, sizeof(magic), 1, f) == 1 + && fwrite(&cnt, sizeof(cnt), 1, f) == 1 + && fwrite(scale_inv, sizeof(float), (size_t)n, f) == (size_t)n; + fclose(f); + return ok; +} + +// Reads scale_inv into out[0..n). Returns the count read, or -1 on error/mismatch. +static inline int turbo_innerq_load_file(const char * path, float * out, int n) { + if (!path || n <= 0 || n > INNERQ_MAX_CHANNELS) return -1; + FILE * f = fopen(path, "rb"); + if (!f) return -1; + uint32_t magic = 0; int32_t cnt = 0; + if (fread(&magic, sizeof(magic), 1, f) != 1 || magic != TURBO_INNERQ_MAGIC || + fread(&cnt, sizeof(cnt), 1, f) != 1 || cnt <= 0 || cnt > n) { fclose(f); return -1; } + const size_t got = fread(out, sizeof(float), (size_t)cnt, f); + fclose(f); + return got == (size_t)cnt ? (int)cnt : -1; +} + +// ---- v2: per-width slots (head_dim varies per layer, e.g. Gemma-4 256-SWA + 512-global) ---- +// Layout: u32 magic(TIQ2), i32 nslots, then per slot { i32 n_embd_gqa, i32 head_dim, +// float[head_dim] scale_inv }. n_embd_gqa is the set-rows row width (ne00) that selects the +// slot at runtime; 0 = wildcard (matches any width — used by the v1→v2 promotion below). +static inline bool turbo_innerq_dump_file_v2(const char * path, int nslots, + const int * nembd, const int * hd, const float scale_inv[][INNERQ_MAX_CHANNELS]) { + if (!path || nslots <= 0 || nslots > INNERQ_MAX_SLOTS) return false; + FILE * f = fopen(path, "wb"); + if (!f) return false; + const uint32_t magic = TURBO_INNERQ_MAGIC2; + const int32_t ns = nslots; + bool ok = fwrite(&magic, sizeof(magic), 1, f) == 1 + && fwrite(&ns, sizeof(ns), 1, f) == 1; + for (int s = 0; s < nslots && ok; s++) { + const int32_t ne = nembd[s], h = hd[s]; + ok = h > 0 && h <= INNERQ_MAX_CHANNELS + && fwrite(&ne, sizeof(ne), 1, f) == 1 + && fwrite(&h, sizeof(h), 1, f) == 1 + && fwrite(scale_inv[s], sizeof(float), (size_t)h, f) == (size_t)h; + } + fclose(f); + return ok; +} + +// Loads per-slot scale_inv. Accepts both v2 (multi-slot) and v1 (single uniform → 1 wildcard +// slot, n_embd_gqa=0). Returns nslots read, or -1 on error. out[s][0..hd_out[s]). +static inline int turbo_innerq_load_file_v2(const char * path, int * nembd_out, int * hd_out, + float out[][INNERQ_MAX_CHANNELS]) { + if (!path) return -1; + FILE * f = fopen(path, "rb"); + if (!f) return -1; + uint32_t magic = 0; + if (fread(&magic, sizeof(magic), 1, f) != 1) { fclose(f); return -1; } + if (magic == TURBO_INNERQ_MAGIC2) { + int32_t ns = 0; + if (fread(&ns, sizeof(ns), 1, f) != 1 || ns <= 0 || ns > INNERQ_MAX_SLOTS) { fclose(f); return -1; } + for (int s = 0; s < ns; s++) { + int32_t ne = 0, h = 0; + if (fread(&ne, sizeof(ne), 1, f) != 1 || fread(&h, sizeof(h), 1, f) != 1 || + h <= 0 || h > INNERQ_MAX_CHANNELS || + fread(out[s], sizeof(float), (size_t)h, f) != (size_t)h) { fclose(f); return -1; } + nembd_out[s] = ne; hd_out[s] = h; + } + fclose(f); + return ns; + } + if (magic == TURBO_INNERQ_MAGIC) { // v1 single uniform → one wildcard slot + int32_t cnt = 0; + if (fread(&cnt, sizeof(cnt), 1, f) != 1 || cnt <= 0 || cnt > INNERQ_MAX_CHANNELS || + fread(out[0], sizeof(float), (size_t)cnt, f) != (size_t)cnt) { fclose(f); return -1; } + nembd_out[0] = 0; hd_out[0] = cnt; + fclose(f); + return 1; + } + fclose(f); + return -1; +} diff --git a/llama.cpp/ggml/src/ggml-cuda/turbo-wht-core.cuh b/llama.cpp/ggml/src/ggml-cuda/turbo-wht-core.cuh new file mode 100644 --- /dev/null +++ b/llama.cpp/ggml/src/ggml-cuda/turbo-wht-core.cuh @@ -0,0 +1,67 @@ +// opencoti-hook: turboquant perf-lift (Level B) — see docs/features/poly_kv.md (F5 M6-S2 #391) +// +// Pure-CUDA core for GGML_OP_TURBO_WHT: sign tables + the cooperative 128-pt Walsh–Hadamard +// group rotation. Deliberately depends on NOTHING but CUDA built-ins (no ggml / common.cuh), +// so the standalone L1 unit (.opencoti/l1-turbo-wht-unit.cu) compiles the IDENTICAL device +// math the shipped kernel (turbo-wht.cu) runs — the round-trip / Parseval proof exercises the +// real recipe, not a paraphrase. +// +// TURBO_WHT_S1/S2 and 1/√128 are byte-identical to TURBO3_WHT_SIGNS1/2 in turbo-dequant.cuh +// (the convention the turbo K/V cache was quantised with). VERIFIED equal at build authoring +// by .opencoti/l1-turbo-wht-unit.cu's table-diff against turbo-dequant.cuh. +#pragma once + +#define TURBO_WHT_INV_SQRT128 0.08838834764831845f // 1/sqrt(128) + +static __constant__ float TURBO_WHT_S1[128] = { + -1.0f, 1.0f, 1.0f, -1.0f, -1.0f, 1.0f, -1.0f, 1.0f, -1.0f, -1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, + 1.0f, -1.0f, 1.0f, -1.0f, 1.0f, -1.0f, -1.0f, 1.0f, 1.0f, 1.0f, -1.0f, 1.0f, 1.0f, -1.0f, -1.0f, -1.0f, + -1.0f, 1.0f, 1.0f, -1.0f, 1.0f, 1.0f, -1.0f, 1.0f, -1.0f, 1.0f, 1.0f, -1.0f, -1.0f, 1.0f, -1.0f, 1.0f, + 1.0f, 1.0f, 1.0f, -1.0f, -1.0f, -1.0f, -1.0f, -1.0f, 1.0f, -1.0f, 1.0f, 1.0f, 1.0f, 1.0f, -1.0f, 1.0f, + -1.0f, -1.0f, 1.0f, -1.0f, -1.0f, -1.0f, 1.0f, -1.0f, -1.0f, -1.0f, 1.0f, -1.0f, -1.0f, -1.0f, 1.0f, 1.0f, + 1.0f, -1.0f, -1.0f, 1.0f, 1.0f, 1.0f, -1.0f, -1.0f, 1.0f, 1.0f, -1.0f, 1.0f, 1.0f, -1.0f, 1.0f, -1.0f, + -1.0f, 1.0f, 1.0f, -1.0f, 1.0f, -1.0f, 1.0f, -1.0f, 1.0f, 1.0f, 1.0f, 1.0f, -1.0f, 1.0f, -1.0f, 1.0f, + 1.0f, -1.0f, 1.0f, 1.0f, -1.0f, -1.0f, -1.0f, -1.0f, -1.0f, 1.0f, 1.0f, -1.0f, 1.0f, 1.0f, -1.0f, 1.0f +}; + +static __constant__ float TURBO_WHT_S2[128] = { + 1.0f, 1.0f, 1.0f, 1.0f, -1.0f, 1.0f, 1.0f, -1.0f, 1.0f, -1.0f, -1.0f, -1.0f, 1.0f, -1.0f, -1.0f, -1.0f, + 1.0f, 1.0f, -1.0f, -1.0f, 1.0f, -1.0f, 1.0f, -1.0f, 1.0f, -1.0f, -1.0f, 1.0f, -1.0f, 1.0f, 1.0f, 1.0f, + 1.0f, 1.0f, -1.0f, -1.0f, -1.0f, 1.0f, -1.0f, -1.0f, -1.0f, -1.0f, -1.0f, -1.0f, 1.0f, 1.0f, 1.0f, -1.0f, + 1.0f, -1.0f, 1.0f, 1.0f, 1.0f, -1.0f, -1.0f, 1.0f, -1.0f, -1.0f, -1.0f, -1.0f, -1.0f, -1.0f, 1.0f, 1.0f, + 1.0f, -1.0f, 1.0f, -1.0f, -1.0f, -1.0f, -1.0f, 1.0f, -1.0f, 1.0f, -1.0f, 1.0f, -1.0f, -1.0f, 1.0f, 1.0f, + -1.0f, 1.0f, -1.0f, 1.0f, 1.0f, -1.0f, 1.0f, -1.0f, -1.0f, -1.0f, -1.0f, 1.0f, -1.0f, -1.0f, 1.0f, -1.0f, + 1.0f, -1.0f, 1.0f, 1.0f, 1.0f, -1.0f, -1.0f, 1.0f, -1.0f, 1.0f, -1.0f, 1.0f, 1.0f, -1.0f, -1.0f, 1.0f, + -1.0f, 1.0f, -1.0f, 1.0f, 1.0f, -1.0f, 1.0f, -1.0f, 1.0f, -1.0f, -1.0f, -1.0f, -1.0f, -1.0f, 1.0f, -1.0f +}; + +// Cooperative 128-pt WHT over shared buf[0..127]; thread t owns buf[t]; all 128 threads +// participate (no divergent early-return before the syncs). Returns lane t's rotated value; +// the caller writes it out. direction 0 = forward (InnerQ ×scale_inv BEFORE signs1→WHT→ +// 1/√128·signs2); direction 1 = inverse (signs2→WHT→1/√128·signs1, InnerQ ×scale_inv AFTER). +// `ch` = InnerQ channel index for this lane (channel within head_dim); scale_inv may be +// nullptr (pure orthonormal WHT). Butterfly is the turbo-cpy.cu validated `(t & h)==0` form. +static __device__ __forceinline__ float turbo_wht_group128( + float * buf, int t, int direction, const float * scale_inv, int ch) { + const float * s_first = (direction == 0) ? TURBO_WHT_S1 : TURBO_WHT_S2; + const float * s_second = (direction == 0) ? TURBO_WHT_S2 : TURBO_WHT_S1; + + __syncthreads(); + if (direction == 0 && scale_inv) buf[t] *= scale_inv[ch]; + __syncthreads(); + buf[t] *= s_first[t]; + + for (int h = 1; h < 128; h <<= 1) { + __syncthreads(); + if ((t & h) == 0) { + const float a = buf[t], b = buf[t + h]; + buf[t] = a + b; + buf[t + h] = a - b; + } + } + __syncthreads(); + + float val = buf[t] * TURBO_WHT_INV_SQRT128 * s_second[t]; + if (direction == 1 && scale_inv) val *= scale_inv[ch]; + return val; +} diff --git a/llama.cpp/ggml/src/ggml-cuda/turbo-wht.cu b/llama.cpp/ggml/src/ggml-cuda/turbo-wht.cu new file mode 100644 --- /dev/null +++ b/llama.cpp/ggml/src/ggml-cuda/turbo-wht.cu @@ -0,0 +1,73 @@ +// opencoti-hook: turboquant perf-lift (Level B) — see docs/features/poly_kv.md (F5 M6-S2 #391) +// +// GGML_OP_TURBO_WHT CUDA forward. One threadblock (128 threads) per 128-element WHT group. +// Rotates the small n_q-sized Q (direction 0) / FA-output (direction 1) tensors — NOT the KV +// cache (that stays compressed and is read in-register by the fused FA-VEC kernel at L2). The +// device math lives in turbo-wht-core.cuh, shared verbatim with the standalone L1 unit so the +// round-trip/Parseval proof exercises the real recipe. +// +// InnerQ is B-adapted: scale_inv comes from the device symbol (turbo_innerq_active_device_scale, +// keyed on ne[0]=head_dim) unless a non-NULL src[1] (fork-compat) overrides it. Default-off / +// width-miss ⇒ nullptr ⇒ pure orthonormal WHT, exactly as the Level-A dequant lift behaves. + +#include "common.cuh" +#include "turbo-wht.cuh" +#include "turbo-innerq-dev.cuh" + +static __global__ void k_turbo_wht_f32( + const float * __restrict__ src, float * __restrict__ dst, int direction, + int64_t head_dim, int64_t groups_per_head, int64_t n_groups, + const float * __restrict__ scale_inv) { + const int64_t g = (int64_t) blockIdx.x; + if (g >= n_groups) return; + const int t = threadIdx.x; + + const int64_t head_idx = g / groups_per_head; + const int64_t grp_in_head = g % groups_per_head; + const int64_t base = head_idx * head_dim + grp_in_head * 128; + + __shared__ float buf[128]; + buf[t] = src[base + t]; + + // InnerQ channel = position within head_dim = grp_in_head*128 + t (== base % head_dim + t). + const int ch = (int) (grp_in_head * 128 + t); + const float r = turbo_wht_group128(buf, t, direction, scale_inv, ch); + dst[base + t] = r; +} + +void ggml_cuda_turbo_wht(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { + const ggml_tensor * src = dst->src[0]; + GGML_ASSERT(src->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32); + GGML_ASSERT(ggml_is_contiguous(src) && ggml_is_contiguous(dst)); + + int direction = 0, group_size = 0; + memcpy(&direction, (const char *) dst->op_params + 0, sizeof(int)); + memcpy(&group_size, (const char *) dst->op_params + sizeof(int), sizeof(int)); + GGML_ASSERT(group_size == 128 && "CUDA turbo_wht supports group_size 128 only (the WHT group)"); + + const int64_t head_dim = src->ne[0]; + GGML_ASSERT(head_dim % 128 == 0); + const int64_t n_elem = ggml_nelements(src); + const int64_t n_heads = n_elem / head_dim; + const int64_t groups_per_head = head_dim / 128; + const int64_t n_groups = groups_per_head * n_heads; + GGML_ASSERT(n_groups > 0 && n_groups < UINT_MAX); + + // B-adapted InnerQ: explicit src[1] graph tensor (fork-compat) wins; else the device symbol. + const float * d_scale_inv = nullptr; + if (dst->src[1]) { + d_scale_inv = (const float *) dst->src[1]->data; + } else { + const float * p = nullptr; + int hd = 0; + if (turbo_innerq_active_device_scale(head_dim, &p, &hd) && hd == (int) head_dim) { + d_scale_inv = p; + } + } + + const float * src_d = (const float *) src->data; + float * dst_d = (float *) dst->data; + cudaStream_t stream = ctx.stream(); + k_turbo_wht_f32<<<(unsigned) n_groups, 128, 0, stream>>>( + src_d, dst_d, direction, head_dim, groups_per_head, n_groups, d_scale_inv); +} diff --git a/llama.cpp/ggml/src/ggml-cuda/turbo-wht.cuh b/llama.cpp/ggml/src/ggml-cuda/turbo-wht.cuh new file mode 100644 --- /dev/null +++ b/llama.cpp/ggml/src/ggml-cuda/turbo-wht.cuh @@ -0,0 +1,27 @@ +// opencoti-hook: turboquant perf-lift (Level B) — see docs/features/poly_kv.md (F5 M6-S2 #391) +// +// GGML_OP_TURBO_WHT: a Walsh–Hadamard rotation op over f32, for the FUSED turbo FA-VEC +// path (Level B). It runs on the small n_q-sized tensors — never on the KV cache: +// direction 0 (forward, Q pre-rotation): [×scale_inv] → signs1 → WHT → 1/√128·signs2 +// direction 1 (inverse, FA-output un-rotation): signs2 → WHT → 1/√128·signs1 → [×scale_inv] +// +// The sign tables (TURBO3_WHT_SIGNS1/2) and 1/√128 normalisation are EXACTLY the +// convention the K/V cache was quantised with (turbo-dequant.cuh / set-rows.cu), so the +// Parseval identity holds: forward-WHT(Q·scale_inv) · stored-rotated-K = Q_true·K_true. +// +// InnerQ is B-adapted (vs the AtomicBot fork's graph-tensor src[1]): the launcher +// resolves scale_inv from the device symbol (turbo_innerq_active_device_scale, keyed on +// ne[0]=head_dim); a non-NULL src[1] (fork-compat) takes precedence. Group size is 128 +// (the WHT group; Gemma head_dim 256/512 = 2/4 groups). Acceptance = logit-equivalence. +#pragma once + +#include "common.cuh" + +// Cooperative in-shared WHT over one 128-element group. All 128 threads of the block +// participate (thread t owns buf[t]); no divergent early-return before the syncs. `ch` is +// the InnerQ channel index for this lane (= channel within head_dim); scale_inv may be +// nullptr (⇒ no InnerQ, pure orthonormal WHT — used by the round-trip unit & InnerQ-off). +// Defined in turbo-wht-core.cuh so the shipped kernel and the standalone unit share it. +#include "turbo-wht-core.cuh" + +void ggml_cuda_turbo_wht(ggml_backend_cuda_context & ctx, ggml_tensor * dst); diff --git a/llama.cpp/ggml/src/ggml-quants.h b/llama.cpp/ggml/src/ggml-quants.h --- a/llama.cpp/ggml/src/ggml-quants.h +++ b/llama.cpp/ggml/src/ggml-quants.h @@ -107,6 +107,31 @@ GGML_API void iq2xs_free_impl(enum ggml_type type); GGML_API void iq3xs_init_impl(int grid_size); GGML_API void iq3xs_free_impl(int grid_size); +// opencoti-hook: turboquant — see docs/features/poly_kv.md (M6-S2) +// TurboQuant KV cache compression (arXiv 2504.19874) +GGML_API void quantize_row_turbo3_0_ref(const float * GGML_RESTRICT x, block_turbo3_0 * GGML_RESTRICT y, int64_t k); +GGML_API void quantize_row_turbo4_0_ref(const float * GGML_RESTRICT x, block_turbo4_0 * GGML_RESTRICT y, int64_t k); +GGML_API void dequantize_row_turbo3_0(const block_turbo3_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API void dequantize_row_turbo4_0(const block_turbo4_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API size_t quantize_turbo3_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_turbo4_0(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_0_ref(const float * GGML_RESTRICT x, block_turbo2_0 * GGML_RESTRICT y, int64_t k); +GGML_API void dequantize_row_turbo2_0(const block_turbo2_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API size_t quantize_turbo2_0(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_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); + +// 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); +GGML_API void dequantize_row_tq3_1s(const block_tq3_1s * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API size_t quantize_tq3_1s(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); + +// TQ4_1S: WHT-rotated 4-bit weight quantization (16-level Lloyd-Max) +GGML_API void quantize_row_tq4_1s_ref(const float * GGML_RESTRICT x, block_tq4_1s * GGML_RESTRICT y, int64_t k); +GGML_API void dequantize_row_tq4_1s(const block_tq4_1s * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API size_t quantize_tq4_1s(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); + #ifdef __cplusplus } #endif diff --git a/llama.cpp/ggml/src/ggml-turbo-quant.c b/llama.cpp/ggml/src/ggml-turbo-quant.c new file mode 100644 --- /dev/null +++ b/llama.cpp/ggml/src/ggml-turbo-quant.c @@ -0,0 +1,1235 @@ +/* + * TurboQuant: KV cache compression via PolarQuant + QJL + * Based on: arXiv 2504.19874 (ICLR 2026) + * + * Implements GGML_TYPE_TURBO2_0 (2-bit), GGML_TYPE_TURBO3_0 (3-bit) and + * GGML_TYPE_TURBO4_0 (4-bit) for use as --cache-type-k turboN in llama-server. + */ + +#include "ggml-quants.h" +#include "ggml-common.h" +#include "ggml-impl.h" + +#define _USE_MATH_DEFINES +#include +#include +#include +#include +#include +#include + +#ifndef M_PI +#define M_PI 3.14159265358979323846 +#endif + +/* Global: WHT group size for CPU quantize path (set by CPU SET_ROWS handler) */ +GGML_API int turbo3_cpu_wht_group_size = 0; + +/* ---------- constants ---------- */ + +#define TURBO_SEED_ROTATION 42 +#define TURBO_SEED_QJL 1042 +#define TURBO_D 128 /* rotation group size = head_dim (independent of block size) */ +#define TURBO_QJL_CONST 1.2533141373155003f /* sqrt(pi/2) */ + +/* Optimal centroids from paper (scaled by 1/sqrt(d)) */ +/* 2-bit: {±0.453, ±1.51} / sqrt(d) */ +static const float CENTROIDS_2BIT[4] = { -0.133462f, -0.039994f, 0.039994f, 0.133462f }; + +/* 3-bit: Lloyd-Max for N(0, 1/128), pre-computed */ +static const float CENTROIDS_3BIT[8] = { + -0.190685f, -0.117832f, -0.065717f, -0.021460f, + 0.021460f, 0.065717f, 0.117832f, 0.190685f +}; + +/* ---------- rotation matrix (lazy init) ---------- */ + +static float turbo_rotation[TURBO_D * TURBO_D]; +static float turbo_rotation_t[TURBO_D * TURBO_D]; /* transpose */ +static int turbo_rotation_initialized = 0; + +/* Simple LCG PRNG for deterministic rotation generation */ +static uint64_t turbo_prng_state; + +static void turbo_prng_seed(uint64_t seed) { + turbo_prng_state = seed; +} + +static double turbo_prng_normal(void) { + /* Box-Muller transform from uniform LCG */ + turbo_prng_state = turbo_prng_state * 6364136223846793005ULL + 1442695040888963407ULL; + double u1 = (double)(turbo_prng_state >> 11) / (double)(1ULL << 53); + if (u1 < 1e-15) u1 = 1e-15; + turbo_prng_state = turbo_prng_state * 6364136223846793005ULL + 1442695040888963407ULL; + double u2 = (double)(turbo_prng_state >> 11) / (double)(1ULL << 53); + return sqrt(-2.0 * log(u1)) * cos(2.0 * M_PI * u2); +} + +static void turbo_init_rotation(void) { + if (turbo_rotation_initialized) return; + + const int d = TURBO_D; + + /* Generate random Gaussian matrix */ + turbo_prng_seed(TURBO_SEED_ROTATION); + float G[TURBO_D * TURBO_D]; + for (int i = 0; i < d * d; i++) { + G[i] = (float)turbo_prng_normal(); + } + + /* QR decomposition via modified Gram-Schmidt */ + /* Q stored column-major in turbo_rotation */ + memcpy(turbo_rotation, G, d * d * sizeof(float)); + + for (int j = 0; j < d; j++) { + /* Normalize column j */ + float norm = 0.0f; + for (int i = 0; i < d; i++) { + norm += turbo_rotation[i * d + j] * turbo_rotation[i * d + j]; + } + norm = sqrtf(norm); + if (norm > 1e-10f) { + for (int i = 0; i < d; i++) { + turbo_rotation[i * d + j] /= norm; + } + } + + /* Orthogonalize remaining columns against j */ + for (int k = j + 1; k < d; k++) { + float dot = 0.0f; + for (int i = 0; i < d; i++) { + dot += turbo_rotation[i * d + j] * turbo_rotation[i * d + k]; + } + for (int i = 0; i < d; i++) { + turbo_rotation[i * d + k] -= dot * turbo_rotation[i * d + j]; + } + } + } + + /* Compute transpose */ + for (int i = 0; i < d; i++) { + for (int j = 0; j < d; j++) { + turbo_rotation_t[i * d + j] = turbo_rotation[j * d + i]; + } + } + + turbo_rotation_initialized = 1; +} + +/* ---------- QJL projection matrix (lazy init, seed-based) ---------- */ + +static float turbo_qjl_matrix[TURBO_D * TURBO_D]; +static float turbo_qjl_matrix_t[TURBO_D * TURBO_D]; +static int turbo_qjl_initialized = 0; + +static void turbo_init_qjl(void) { + if (turbo_qjl_initialized) return; + + const int d = TURBO_D; + turbo_prng_seed(TURBO_SEED_QJL); + + for (int i = 0; i < d * d; i++) { + turbo_qjl_matrix[i] = (float)turbo_prng_normal(); + } + + /* Transpose */ + for (int i = 0; i < d; i++) { + for (int j = 0; j < d; j++) { + turbo_qjl_matrix_t[i * d + j] = turbo_qjl_matrix[j * d + i]; + } + } + + turbo_qjl_initialized = 1; +} + +/* ---------- helper: matrix-vector multiply ---------- */ + +static void matvec(const float * M, const float * x, float * y, int d) { + /* y = M @ x, M is row-major d×d */ + for (int i = 0; i < d; i++) { + float sum = 0.0f; + for (int j = 0; j < d; j++) { + sum += M[i * d + j] * x[j]; + } + y[i] = sum; + } +} + +/* ---------- nearest centroid ---------- */ + +static int nearest_centroid_2bit(float val) { + /* Binary search on midpoints: {-0.133, -0.040, 0.040, 0.133} */ + if (val < -0.086728f) return 0; /* midpoint(-0.133, -0.040) */ + if (val < 0.000000f) return 1; /* midpoint(-0.040, 0.040) */ + if (val < 0.086728f) return 2; /* midpoint(0.040, 0.133) */ + return 3; +} + +static int nearest_centroid_3bit(float val) { + /* 8 centroids, find nearest via midpoints */ + if (val < -0.154259f) return 0; + if (val < -0.091775f) return 1; + if (val < -0.043589f) return 2; + if (val < 0.000000f) return 3; + if (val < 0.043589f) return 4; + if (val < 0.091775f) return 5; + if (val < 0.154259f) return 6; + return 7; +} + +static int nearest_centroid_4bit(float val) { + /* 16 centroids, optimal for N(0, 1/sqrt(128)), find nearest via midpoints */ + if (val < -0.145560f) return 0; + if (val < -0.103361f) return 1; + if (val < -0.079142f) return 2; + if (val < -0.060009f) return 3; + if (val < -0.043430f) return 4; + if (val < -0.028293f) return 5; + if (val < -0.013963f) return 6; + if (val < 0.000000f) return 7; + if (val < 0.013963f) return 8; + if (val < 0.028293f) return 9; + if (val < 0.043430f) return 10; + if (val < 0.060009f) return 11; + if (val < 0.079142f) return 12; + if (val < 0.103361f) return 13; + if (val < 0.145560f) return 14; + return 15; +} + +/* ---------- WHT sign arrays (must match CUDA/Metal, seed=42) ---------- */ + +static const float turbo_cpu_s1[128] = { + -1,1,1,-1,-1,1,-1,1,-1,-1,1,1,1,1,1,1,1,-1,1,-1,1,-1,-1,1,1,1,-1,1,1,-1,-1,-1, + -1,1,1,-1,1,1,-1,1,-1,1,1,-1,-1,1,-1,1,1,1,1,-1,-1,-1,-1,-1,1,-1,1,1,1,1,-1,1, + -1,-1,1,-1,-1,-1,1,-1,-1,-1,1,-1,-1,-1,1,1,1,-1,-1,1,1,1,-1,-1,1,1,-1,1,1,-1,1,-1, + -1,1,1,-1,1,-1,1,-1,1,1,1,1,-1,1,-1,1,1,-1,1,1,-1,-1,-1,-1,-1,1,1,-1,1,1,-1,1 +}; + +static const float turbo_cpu_s2[128] = { + 1,1,1,1,-1,1,1,-1,1,-1,-1,-1,1,-1,-1,-1,1,1,-1,-1,1,-1,1,-1,1,-1,-1,1,-1,1,1,1, + 1,1,-1,-1,-1,1,-1,-1,-1,-1,-1,-1,1,1,1,-1,1,-1,1,1,1,-1,-1,1,-1,-1,-1,-1,-1,-1,1,1, + 1,-1,1,-1,-1,-1,-1,1,-1,1,-1,1,-1,-1,1,1,-1,1,-1,1,1,-1,1,-1,-1,-1,-1,1,-1,-1,1,-1, + 1,-1,1,1,1,-1,-1,1,-1,1,-1,1,1,-1,-1,1,-1,1,-1,1,1,-1,1,-1,1,-1,-1,-1,-1,-1,1,-1 +}; + +/* ---------- CPU forward WHT (in-place, group_size elements) ---------- */ + +static void turbo_cpu_fwht(float * x, int group_size) { + const float * s1 = turbo_cpu_s1; + const float * s2 = turbo_cpu_s2; + const float inv_sqrt = (group_size == 128) ? 0.08838834764831845f : 0.125f; + + // signs1 + for (int i = 0; i < group_size; i++) x[i] *= s1[i]; + + // butterfly stages + for (int h = 1; h < group_size; h *= 2) { + for (int i = 0; i < group_size; i += h * 2) { + for (int j = i; j < i + h; j++) { + float a = x[j], b = x[j + h]; + x[j] = a + b; + x[j + h] = a - b; + } + } + } + + // normalize + signs2 + for (int i = 0; i < group_size; i++) x[i] *= inv_sqrt * s2[i]; +} + +/* ---------- CPU inverse WHT (in-place, group_size elements) ---------- + * + * Forward is y = D(s2) * N * H * D(s1) * x (N = 1/sqrt(group_size)) + * H is the unnormalized Hadamard butterfly with H*H = group_size * I, so + * (N*H) is self-inverse. s1 and s2 are ±1 diagonals, also self-inverse. + * The inverse therefore has the same structure with s1 and s2 swapped: + * x = D(s1) * N * H * D(s2) * y + */ +GGML_API void turbo_cpu_fwht_inverse(float * x, int group_size) { + const float * s1 = turbo_cpu_s1; + const float * s2 = turbo_cpu_s2; + const float inv_sqrt = (group_size == 128) ? 0.08838834764831845f : 0.125f; + + // signs2 (undoes the s2 that was applied last in the forward pass) + for (int i = 0; i < group_size; i++) x[i] *= s2[i]; + + // butterfly stages (same as forward — self-inverse up to the inv_sqrt scaling below) + for (int h = 1; h < group_size; h *= 2) { + for (int i = 0; i < group_size; i += h * 2) { + for (int j = i; j < i + h; j++) { + float a = x[j], b = x[j + h]; + x[j] = a + b; + x[j + h] = a - b; + } + } + } + + // normalize + signs1 + for (int i = 0; i < group_size; i++) x[i] *= inv_sqrt * s1[i]; +} + +/* ---------- TURBO3_0: 3-bit PolarQuant with WHT rotation ---------- */ + +void quantize_row_turbo3_0_ref(const float * GGML_RESTRICT x, block_turbo3_0 * GGML_RESTRICT y, int64_t k) { + assert(k % QK_TURBO3 == 0); + + // Read WHT group size from global (set by CPU SET_ROWS handler before each call). + // Fallback: 128 if row is 128-aligned, else 64. + extern int turbo3_cpu_wht_group_size; + int group_size = turbo3_cpu_wht_group_size; + if (group_size != 64 && group_size != 128) { + group_size = (k % 128 == 0) ? 128 : 64; + } + if (k % group_size != 0) group_size = (group_size == 128) ? 64 : 128; + assert(k % group_size == 0); + + const int n_groups = k / group_size; + const int blocks_per_group = group_size / QK_TURBO3; + + for (int g = 0; g < n_groups; g++) { + const float * grp_src = x + g * group_size; + block_turbo3_0 * grp_dst = y + g * blocks_per_group; + + // 1. L2 norm over the group + float norm_sq = 0.0f; + float buf[128]; // max group_size + for (int j = 0; j < group_size; j++) { + buf[j] = grp_src[j]; + norm_sq += buf[j] * buf[j]; + } + float grp_norm = sqrtf(norm_sq); + float inv_norm = (grp_norm > 1e-10f) ? 1.0f / grp_norm : 0.0f; + + // 2. Normalize + for (int j = 0; j < group_size; j++) buf[j] *= inv_norm; + + // 3. Forward WHT rotation + turbo_cpu_fwht(buf, group_size); + + // 4. Quantize + pack into sub-blocks + float recon_sq = 0.0f; + for (int b = 0; b < blocks_per_group; b++) { + block_turbo3_0 * blk = &grp_dst[b]; + const int off = b * QK_TURBO3; + + memset(blk->qs, 0, QK_TURBO3 / 4); + memset(blk->signs, 0, QK_TURBO3 / 8); + + for (int j = 0; j < QK_TURBO3; j++) { + int idx = nearest_centroid_3bit(buf[off + j]); + blk->qs[j / 4] |= (idx & 0x3) << ((j % 4) * 2); + if (idx & 0x4) { + blk->signs[j / 8] |= (1 << (j % 8)); + } + recon_sq += CENTROIDS_3BIT[idx] * CENTROIDS_3BIT[idx]; + } + } + + // 5. Corrected norm: grp_norm / recon_norm (matching CUDA kernel) + float recon_norm = sqrtf(recon_sq); + float corrected = (recon_norm > 1e-10f) ? grp_norm / recon_norm : grp_norm; + for (int b = 0; b < blocks_per_group; b++) { + grp_dst[b].norm = GGML_FP32_TO_FP16(corrected); + } + } +} + +/* opencoti-hook: turboquant InnerQ (M6-S2) — CPU dequant per-channel un-equalize, + * PER-WIDTH (per-layer head_dim). Loads the TURBO_INNERQ_SCALE artifact once: v2 + * (magic "TIQ2" = 0x32514954) carries one slot per distinct KV width (n_embd_gqa → + * head_dim → scale_inv[head_dim]); v1 (magic "TIQ1" = 0x31514954) is a single uniform + * vector promoted to one wildcard slot (n_embd_gqa=0). Identity if unset/unreadable ⇒ + * default-off is byte-identical. The GPU set-rows forward equalizes K (path-shared); + * THIS un-equalize is the dequant-on-lift half of the bridge. + * + * FUSED-PATH INVARIANT (next milestone): this ×scale_inv un-equalize must stay isolated + * to dequantize_row_turbo3_0 (the lift's to_float). The fused attention path keeps K + * equalized and instead pre-scales Q by scale_inv (Q_pre·K_eq = Q·K), reading the SAME + * per-width table — so it must NOT route through this to_float (no double-correction). + * Do not move ×scale_inv into any code shared with the fused kernel. */ +#define INNERQ_CPU_MAX 512 /* max head_dim (Gemma global = 512) */ +#define INNERQ_CPU_MAX_SLOTS 4 /* distinct KV widths (Gemma-4 = 2: SWA+global) */ +static int g_innerq_cpu_nslots = 0; +static int g_innerq_cpu_nembd[INNERQ_CPU_MAX_SLOTS]; +static int g_innerq_cpu_hd [INNERQ_CPU_MAX_SLOTS]; +static float g_innerq_cpu_sinv [INNERQ_CPU_MAX_SLOTS][INNERQ_CPU_MAX]; +static void turbo_innerq_cpu_load(void) { + static int loaded = 0; + if (loaded) return; + int ns = 0; + const char * path = getenv("TURBO_INNERQ_SCALE"); + if (path && path[0]) { + FILE * f = fopen(path, "rb"); + if (f) { + uint32_t magic = 0; + if (fread(&magic, sizeof(magic), 1, f) == 1) { + if (magic == 0x32514954u) { /* TIQ2: per-width slots */ + int32_t nslots = 0; + if (fread(&nslots, sizeof(nslots), 1, f) == 1 && + nslots > 0 && nslots <= INNERQ_CPU_MAX_SLOTS) { + int ok = 1; + for (int s = 0; s < nslots && ok; s++) { + int32_t ne = 0, hd = 0; + if (fread(&ne, sizeof(ne), 1, f) == 1 && fread(&hd, sizeof(hd), 1, f) == 1 && + hd > 0 && hd <= INNERQ_CPU_MAX && + fread(g_innerq_cpu_sinv[s], sizeof(float), (size_t)hd, f) == (size_t)hd) { + g_innerq_cpu_nembd[s] = ne; g_innerq_cpu_hd[s] = hd; + } else ok = 0; + } + if (ok) ns = nslots; + } + } else if (magic == 0x31514954u) { /* TIQ1: single uniform → wildcard slot */ + int32_t cnt = 0; + if (fread(&cnt, sizeof(cnt), 1, f) == 1 && cnt > 0 && cnt <= INNERQ_CPU_MAX && + fread(g_innerq_cpu_sinv[0], sizeof(float), (size_t)cnt, f) == (size_t)cnt) { + g_innerq_cpu_nembd[0] = 0; g_innerq_cpu_hd[0] = cnt; ns = 1; + } + } + } + fclose(f); + } + } + g_innerq_cpu_nslots = ns; /* 0 ⇒ identity (no InnerQ) */ + loaded = 1; /* publish only after the table is fully populated */ +} +/* Select the slot for a row of length k (= n_embd_gqa): exact width match, then a + * wildcard (n_embd_gqa==0) slot, else -1 (no InnerQ for this width). */ +static int turbo_innerq_cpu_slot_for(int64_t k) { + for (int s = 0; s < g_innerq_cpu_nslots; s++) if (g_innerq_cpu_nembd[s] == (int)k) return s; + for (int s = 0; s < g_innerq_cpu_nslots; s++) if (g_innerq_cpu_nembd[s] == 0) return s; + return -1; +} + +void dequantize_row_turbo3_0(const block_turbo3_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { + // opencoti-hook: turboquant (M6-S2) — faithful CPU dequant (was a rotated-space + // stub that skipped the inverse WHT). Forward is: + // normalize → forward WHT (turbo_cpu_fwht) → quantize-to-centroid + // so the stored centroids are post-rotation, normalized values. The true-space + // inverse must therefore: reconstruct centroids → inverse WHT → denormalize (×norm). + // The previous `CENTROIDS_3BIT[idx]*norm` left the data WHT-rotated, which is only + // correct for a fused kernel that pre-rotates Q; for a plain dequant→f16→stock-FA + // path (and any offline round-trip) it is wrong. QK_TURBO3==128 is exactly one WHT + // group, so per-block inverse == per-group inverse. + assert(k % QK_TURBO3 == 0); + const int nb = k / QK_TURBO3; + // Per-width InnerQ: k == n_embd_gqa selects this layer's slot (head_dim + scale_inv). + // slot<0 ⇒ no InnerQ for this width (or default-off) ⇒ ×scale_inv is identity. + turbo_innerq_cpu_load(); + const int slot = turbo_innerq_cpu_slot_for(k); + const float * sinv = (slot >= 0) ? g_innerq_cpu_sinv[slot] : NULL; + const int hd = (slot >= 0) ? g_innerq_cpu_hd[slot] : 0; + for (int block = 0; block < nb; block++) { + float norm = GGML_FP16_TO_FP32(x[block].norm); + float buf[QK_TURBO3]; + for (int j = 0; j < QK_TURBO3; j++) { + uint8_t low2 = (x[block].qs[j/4] >> ((j%4)*2)) & 0x3; + uint8_t hi1 = (x[block].signs[j/8] >> (j%8)) & 0x1; + uint8_t idx = low2 | (hi1 << 2); + buf[j] = CENTROIDS_3BIT[idx]; + } + turbo_cpu_fwht_inverse(buf, QK_TURBO3); + for (int j = 0; j < QK_TURBO3; j++) { + // head_dim-aware channel: the dequant processes 128-blocks; folding the global + // element index by this layer's head_dim recovers the per-head channel (scale + // shared across heads). ×norm (block scale) then ×scale_inv[ch] (InnerQ) → true K. + const float si = (sinv && hd > 0) ? sinv[(int)(((int64_t)block * QK_TURBO3 + j) % hd)] : 1.0f; + y[block * QK_TURBO3 + j] = buf[j] * norm * si; + } + } +} + +size_t quantize_turbo3_0(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 == 0); + + size_t row_size = (n_per_row / QK_TURBO3) * sizeof(block_turbo3_0); + for (int64_t row = 0; row < nrows; row++) { + quantize_row_turbo3_0_ref( + src + row * n_per_row, + (block_turbo3_0 *)((char *)dst + row * row_size), + n_per_row + ); + } + return nrows * row_size; +} + +/* ---------- TURBO2_0: 2-bit PolarQuant (no QJL) ---------- */ + +void quantize_row_turbo2_0_ref(const float * GGML_RESTRICT x, block_turbo2_0 * GGML_RESTRICT y, int64_t k) { + assert(k % QK_TURBO2 == 0); + + extern int turbo3_cpu_wht_group_size; + int group_size = turbo3_cpu_wht_group_size; + if (group_size != 64 && group_size != 128) { + group_size = (k % 128 == 0) ? 128 : 64; + } + if (k % group_size != 0) group_size = (group_size == 128) ? 64 : 128; + assert(k % group_size == 0); + + const int n_groups = k / group_size; + const int blocks_per_group = group_size / QK_TURBO2; + + for (int g = 0; g < n_groups; g++) { + const float * grp_src = x + g * group_size; + block_turbo2_0 * grp_dst = y + g * blocks_per_group; + + /* 1. L2 norm over the group */ + float norm_sq = 0.0f; + float buf[128]; + for (int j = 0; j < group_size; j++) { + buf[j] = grp_src[j]; + norm_sq += buf[j] * buf[j]; + } + float grp_norm = sqrtf(norm_sq); + float inv_norm = (grp_norm > 1e-10f) ? 1.0f / grp_norm : 0.0f; + + /* 2. Normalize */ + for (int j = 0; j < group_size; j++) buf[j] *= inv_norm; + + /* 3. Forward WHT rotation */ + turbo_cpu_fwht(buf, group_size); + + /* 4. Quantize + pack into sub-blocks */ + float recon_sq = 0.0f; + for (int b = 0; b < blocks_per_group; b++) { + block_turbo2_0 * blk = &grp_dst[b]; + const int off = b * QK_TURBO2; + + memset(blk->qs, 0, QK_TURBO2 / 4); + + for (int j = 0; j < QK_TURBO2; j++) { + int idx = nearest_centroid_2bit(buf[off + j]); + blk->qs[j / 4] |= (idx & 0x3) << ((j % 4) * 2); + recon_sq += CENTROIDS_2BIT[idx] * CENTROIDS_2BIT[idx]; + } + } + + /* 5. Corrected norm */ + float recon_norm = sqrtf(recon_sq); + float corrected = (recon_norm > 1e-10f) ? grp_norm / recon_norm : grp_norm; + for (int b = 0; b < blocks_per_group; b++) { + grp_dst[b].norm = GGML_FP32_TO_FP16(corrected); + } + } +} + +void dequantize_row_turbo2_0(const block_turbo2_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { + // opencoti-hook: turboquant (M6-S2) — faithful 2-bit CPU dequant for the dequant-on-lift + // path (was a rotated-domain stub: CENTROIDS_2BIT[idx]*norm). 2-bit centroid → inverse WHT → + // ×norm → ×scale_inv(InnerQ), mirroring turbo3. QK_TURBO2==128 == one WHT group ⇒ per-block + // inverse == per-group. No QJL signs at 2-bit. Per-width InnerQ via k==n_embd_gqa slot. + assert(k % QK_TURBO2 == 0); + const int nb = k / QK_TURBO2; + turbo_innerq_cpu_load(); + const int slot = turbo_innerq_cpu_slot_for(k); + const float * sinv = (slot >= 0) ? g_innerq_cpu_sinv[slot] : NULL; + const int hd = (slot >= 0) ? g_innerq_cpu_hd[slot] : 0; + for (int block = 0; block < nb; block++) { + float norm = GGML_FP16_TO_FP32(x[block].norm); + float buf[QK_TURBO2]; + for (int j = 0; j < QK_TURBO2; j++) { + uint8_t idx = (x[block].qs[j/4] >> ((j%4)*2)) & 0x3; + buf[j] = CENTROIDS_2BIT[idx]; + } + turbo_cpu_fwht_inverse(buf, QK_TURBO2); + for (int j = 0; j < QK_TURBO2; j++) { + const float si = (sinv && hd > 0) ? sinv[(int)(((int64_t)block * QK_TURBO2 + j) % hd)] : 1.0f; + y[block * QK_TURBO2 + j] = buf[j] * norm * si; + } + } +} + +size_t quantize_turbo2_0(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 == 0); + + size_t row_size = (n_per_row / QK_TURBO2) * sizeof(block_turbo2_0); + for (int64_t row = 0; row < nrows; row++) { + quantize_row_turbo2_0_ref( + src + row * n_per_row, + (block_turbo2_0 *)((char *)dst + row * row_size), + n_per_row + ); + } + return nrows * row_size; +} + +/* ---------- TURBO8_0: 8-bit uniform int8 (WHT-rotated, near-lossless rung) ---------- */ +/* opencoti-hook: turboquant tier family (M6-S2). Same normalize→forward-WHT→corrected-norm + * pipeline as turbo2/3/4, but quantizes each post-WHT value with a uniform signed int8 + * (q = clamp(round(x*TURBO8_SCALE), -127, 127); centroid = q/TURBO8_SCALE). 1 byte/value, no + * codebook, no QJL. block == group == 128 ⇒ per-block inverse WHT == per-group. */ + +void quantize_row_turbo8_0_ref(const float * GGML_RESTRICT x, block_turbo8_0 * GGML_RESTRICT y, int64_t k) { + assert(k % QK_TURBO8 == 0); + + extern int turbo3_cpu_wht_group_size; + int group_size = turbo3_cpu_wht_group_size; + if (group_size != 64 && group_size != 128) { + group_size = (k % 128 == 0) ? 128 : 64; + } + if (k % group_size != 0) group_size = (group_size == 128) ? 64 : 128; + assert(k % group_size == 0); + + const int n_groups = k / group_size; + const int blocks_per_group = group_size / QK_TURBO8; + + for (int g = 0; g < n_groups; g++) { + const float * grp_src = x + g * group_size; + block_turbo8_0 * grp_dst = y + g * blocks_per_group; + + float norm_sq = 0.0f; + float buf[128]; + for (int j = 0; j < group_size; j++) { + buf[j] = grp_src[j]; + norm_sq += buf[j] * buf[j]; + } + float grp_norm = sqrtf(norm_sq); + float inv_norm = (grp_norm > 1e-10f) ? 1.0f / grp_norm : 0.0f; + + for (int j = 0; j < group_size; j++) buf[j] *= inv_norm; + + turbo_cpu_fwht(buf, group_size); + + float recon_sq = 0.0f; + for (int b = 0; b < blocks_per_group; b++) { + block_turbo8_0 * blk = &grp_dst[b]; + const int off = b * QK_TURBO8; + for (int j = 0; j < QK_TURBO8; j++) { + int q = (int) lroundf(buf[off + j] * TURBO8_SCALE); + if (q > 127) q = 127; + if (q < -127) q = -127; + blk->qs[j] = (int8_t) q; + const float c = (float) q / TURBO8_SCALE; + recon_sq += c * c; + } + } + + float recon_norm = sqrtf(recon_sq); + float corrected = (recon_norm > 1e-10f) ? grp_norm / recon_norm : grp_norm; + for (int b = 0; b < blocks_per_group; b++) { + grp_dst[b].norm = GGML_FP32_TO_FP16(corrected); + } + } +} + +void dequantize_row_turbo8_0(const block_turbo8_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { + // int8 centroid (q/TURBO8_SCALE) → inverse WHT → ×norm → ×scale_inv(InnerQ). Mirrors turbo3. + assert(k % QK_TURBO8 == 0); + const int nb = k / QK_TURBO8; + turbo_innerq_cpu_load(); + const int slot = turbo_innerq_cpu_slot_for(k); + const float * sinv = (slot >= 0) ? g_innerq_cpu_sinv[slot] : NULL; + const int hd = (slot >= 0) ? g_innerq_cpu_hd[slot] : 0; + for (int block = 0; block < nb; block++) { + float norm = GGML_FP16_TO_FP32(x[block].norm); + float buf[QK_TURBO8]; + for (int j = 0; j < QK_TURBO8; j++) buf[j] = (float) x[block].qs[j] / TURBO8_SCALE; + turbo_cpu_fwht_inverse(buf, QK_TURBO8); + for (int j = 0; j < QK_TURBO8; j++) { + const float si = (sinv && hd > 0) ? sinv[(int)(((int64_t)block * QK_TURBO8 + j) % hd)] : 1.0f; + y[block * QK_TURBO8 + j] = buf[j] * norm * si; + } + } +} + +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) { + GGML_UNUSED(imatrix); + assert(n_per_row % QK_TURBO8 == 0); + + size_t row_size = (n_per_row / QK_TURBO8) * sizeof(block_turbo8_0); + for (int64_t row = 0; row < nrows; row++) { + quantize_row_turbo8_0_ref( + src + row * n_per_row, + (block_turbo8_0 *)((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) { + turbo_init_rotation(); + turbo_init_qjl(); + + assert(k % QK_TURBO4 == 0); + const int nb = k / QK_TURBO4; + const int d = QK_TURBO4; + + for (int block = 0; block < nb; block++) { + const float * src = x + block * d; + + /* Step 1: Extract norm */ + float norm_sq = 0.0f; + for (int i = 0; i < d; i++) norm_sq += src[i] * src[i]; + float norm = sqrtf(norm_sq); + + /* Normalize */ + float normalized[TURBO_D]; + if (norm > 1e-10f) { + const float inv = 1.0f / norm; + for (int i = 0; i < d; i++) normalized[i] = src[i] * inv; + } else { + memset(normalized, 0, d * sizeof(float)); + } + + /* Step 2: Forward WHT rotation (matches CUDA set_rows) */ + float rotated[TURBO_D]; + memcpy(rotated, normalized, d * sizeof(float)); + turbo_cpu_fwht(rotated, d); + +#if TURBO4_USE_4BIT + /* Step 3: 4-bit quantization (16 centroids) */ + static const float CENTROIDS_4BIT[16] = { + -0.173926f, -0.117195f, -0.089527f, -0.068756f, + -0.051262f, -0.035597f, -0.020989f, -0.006938f, + 0.006938f, 0.020989f, 0.035597f, 0.051262f, + 0.068756f, 0.089527f, 0.117195f, 0.173926f + }; + uint8_t indices[TURBO_D]; + for (int i = 0; i < d; i++) { + indices[i] = (uint8_t)nearest_centroid_4bit(rotated[i]); + } + + /* Norm correction */ + float recon_norm_sq = 0.0f; + for (int i = 0; i < d; i++) { + recon_norm_sq += CENTROIDS_4BIT[indices[i]] * CENTROIDS_4BIT[indices[i]]; + } + float recon_norm = sqrtf(recon_norm_sq); + float corrected_norm = (recon_norm > 1e-10f) ? norm / recon_norm : norm; + y[block].norm = GGML_FP32_TO_FP16(corrected_norm); +#else + /* Step 3: 3-bit quantization (8 centroids) */ + uint8_t indices[TURBO_D]; + for (int i = 0; i < d; i++) { + indices[i] = (uint8_t)nearest_centroid_3bit(rotated[i]); + } + + /* Step 4: Residual */ + float reconstructed[TURBO_D]; + for (int i = 0; i < d; i++) { + reconstructed[i] = CENTROIDS_3BIT[indices[i]]; + } + float mse_recon[TURBO_D]; + matvec(turbo_rotation_t, reconstructed, mse_recon, d); + + float residual[TURBO_D]; + for (int i = 0; i < d; i++) { + residual[i] = normalized[i] - mse_recon[i]; + } + + /* Step 5: QJL */ + float projected[TURBO_D]; + matvec(turbo_qjl_matrix, residual, projected, d); +#endif + + /* Pack */ +#if !TURBO4_USE_4BIT + y[block].norm = GGML_FP32_TO_FP16(norm); +#endif + +#if TURBO4_USE_4BIT + /* 4-bit PolarQuant: nibble pack into qs[64] */ + memset(y[block].qs, 0, d / 2); + for (int i = 0; i < d; i++) { + y[block].qs[i / 2] |= (uint8_t)((indices[i] & 0xF) << ((i % 2) * 4)); + } + y[block].rnorm = GGML_FP32_TO_FP16(0.0f); +#else + /* Legacy 3-bit + QJL: pack 3-bit indices + QJL signs */ + memset(y[block].qs, 0, d * 3 / 8); + for (int i = 0; i < d; i++) { + int bit_offset = i * 3; + int byte_idx = bit_offset / 8; + int bit_pos = bit_offset % 8; + uint16_t val = (uint16_t)(indices[i] & 0x7); + y[block].qs[byte_idx] |= (uint8_t)(val << bit_pos); + if (bit_pos > 5 && byte_idx + 1 < d * 3 / 8) { + y[block].qs[byte_idx + 1] |= (uint8_t)(val >> (8 - bit_pos)); + } + } + memset(y[block].signs, 0, d / 8); + for (int i = 0; i < d; i++) { + if (projected[i] >= 0.0f) { + y[block].signs[i / 8] |= (1 << (i % 8)); + } + } +#endif + } +} + +void dequantize_row_turbo4_0(const block_turbo4_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { + turbo_init_rotation(); + + assert(k % QK_TURBO4 == 0); + const int nb = k / QK_TURBO4; + const int d = QK_TURBO4; + +#if TURBO4_USE_4BIT + /* opencoti-hook: turboquant (M6-S2) — faithful 4-bit CPU dequant for the dequant-on-lift + * path (was a rotated-domain stub: CENTROIDS_4BIT[idx]*norm, valid only for a fused kernel + * that pre-rotates Q). 4-bit centroid → inverse WHT → ×norm → ×scale_inv(InnerQ), exactly + * mirroring dequantize_row_turbo3_0. QK_TURBO4==128 == one WHT group ⇒ per-block inverse == + * per-group inverse. The forward (GPU k_set_rows_turbo4) writes {corrected_norm, nibbles}; + * this reconstructs TRUE K. Per-width InnerQ: k==n_embd_gqa selects the layer's slot. */ + static const float CENTROIDS_4BIT[16] = { + -0.173926f, -0.117195f, -0.089527f, -0.068756f, + -0.051262f, -0.035597f, -0.020989f, -0.006938f, + 0.006938f, 0.020989f, 0.035597f, 0.051262f, + 0.068756f, 0.089527f, 0.117195f, 0.173926f + }; + turbo_innerq_cpu_load(); + const int slot = turbo_innerq_cpu_slot_for(k); + const float * sinv = (slot >= 0) ? g_innerq_cpu_sinv[slot] : NULL; + const int hd = (slot >= 0) ? g_innerq_cpu_hd[slot] : 0; + for (int block = 0; block < nb; block++) { + float norm = GGML_FP16_TO_FP32(x[block].norm); + float buf[QK_TURBO4]; + for (int j = 0; j < d; j++) { + uint8_t idx = (x[block].qs[j / 2] >> ((j % 2) * 4)) & 0xF; + buf[j] = CENTROIDS_4BIT[idx]; + } + turbo_cpu_fwht_inverse(buf, QK_TURBO4); + for (int j = 0; j < d; j++) { + const float si = (sinv && hd > 0) ? sinv[(int)(((int64_t)block * QK_TURBO4 + j) % hd)] : 1.0f; + y[block * d + j] = buf[j] * norm * si; + } + } +#else + /* Legacy 3-bit + QJL dequant */ + turbo_init_qjl(); + for (int block = 0; block < nb; block++) { + float norm = GGML_FP16_TO_FP32(x[block].norm); + + uint8_t indices[TURBO_D]; + for (int i = 0; i < d; i++) { + int bit_offset = i * 3; + int byte_idx = bit_offset / 8; + int bit_pos = bit_offset % 8; + uint16_t raw = (uint16_t)x[block].qs[byte_idx]; + if (byte_idx + 1 < d * 3 / 8) { + raw |= (uint16_t)x[block].qs[byte_idx + 1] << 8; + } + indices[i] = (uint8_t)((raw >> bit_pos) & 0x7); + } + + float signs[TURBO_D]; + for (int i = 0; i < d; i++) { + signs[i] = (x[block].signs[i / 8] & (1 << (i % 8))) ? 1.0f : -1.0f; + } + + float rnorm = GGML_FP16_TO_FP32(x[block].rnorm); + const float qjl_scale = TURBO_QJL_CONST / (float)d * rnorm; + + float rotated_recon[TURBO_D]; + for (int i = 0; i < d; i++) { + rotated_recon[i] = CENTROIDS_3BIT[indices[i]]; + } + float mse_recon[TURBO_D]; + matvec(turbo_rotation_t, rotated_recon, mse_recon, d); + + float qjl_recon[TURBO_D]; + matvec(turbo_qjl_matrix_t, signs, qjl_recon, d); + for (int i = 0; i < d; i++) { + qjl_recon[i] *= qjl_scale; + } + + float * dst = y + block * d; + for (int i = 0; i < d; i++) { + dst[i] = (mse_recon[i] + qjl_recon[i]) * norm; + } + } +#endif +} + +size_t quantize_turbo4_0(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_TURBO4 == 0); + + size_t row_size = (n_per_row / QK_TURBO4) * sizeof(block_turbo4_0); + for (int64_t row = 0; row < nrows; row++) { + quantize_row_turbo4_0_ref( + src + row * n_per_row, + (block_turbo4_0 *)((char *)dst + row * row_size), + n_per_row + ); + } + return nrows * row_size; +} + +/* ================================================================== */ +/* TQ3_1S / TQ4_1S: WHT-rotated weight quantization */ +/* ================================================================== */ + +/* Lloyd-Max centroids for N(0,1) — shared with Metal shaders */ +static const float TQ3_0_CENTROIDS[8] = { + -1.996684f, -1.291398f, -0.740341f, -0.247508f, + 0.230106f, 0.725222f, 1.277503f, 1.988943f +}; + +static const float TQ4_0_CENTROIDS[16] = { + -2.732590f, -2.069017f, -1.618046f, -1.256231f, + -0.942340f, -0.656759f, -0.388048f, -0.128395f, + 0.128395f, 0.388048f, 0.656759f, 0.942340f, + 1.256231f, 1.618046f, 2.069017f, 2.732590f, +}; + +/* WHT sign pattern (golden ratio hash, 32-element blocks) — shared by TQ3 and TQ4 */ +static const float TQ3_0_SIGNS[32] = { + +1.0f, -1.0f, +1.0f, -1.0f, +1.0f, +1.0f, -1.0f, +1.0f, + -1.0f, -1.0f, +1.0f, -1.0f, +1.0f, +1.0f, -1.0f, +1.0f, + -1.0f, -1.0f, +1.0f, -1.0f, +1.0f, -1.0f, -1.0f, +1.0f, + -1.0f, +1.0f, +1.0f, -1.0f, +1.0f, -1.0f, -1.0f, +1.0f, +}; + +#define TQ_BLOCK_SIZE 32 +#define TQ_INV_SQRT32 0.17677669529663688f /* 1/sqrt(32) */ + +/* Forward RHT: sign flips -> WHT butterfly -> normalize */ +static void tq3_0_rht_forward(float * buf) { + for (int i = 0; i < TQ_BLOCK_SIZE; i++) buf[i] *= TQ3_0_SIGNS[i]; + for (int step = 1; step < TQ_BLOCK_SIZE; step <<= 1) { + for (int i = 0; i < TQ_BLOCK_SIZE; i += step << 1) { + for (int j = i; j < i + step; j++) { + float a = buf[j], b = buf[j + step]; + buf[j] = a + b; + buf[j + step] = a - b; + } + } + } + for (int i = 0; i < TQ_BLOCK_SIZE; i++) buf[i] *= TQ_INV_SQRT32; +} + +/* Inverse RHT: WHT butterfly -> normalize + unsign */ +static void tq3_0_rht_inverse(float * buf) { + for (int step = 1; step < TQ_BLOCK_SIZE; step <<= 1) { + for (int i = 0; i < TQ_BLOCK_SIZE; i += step << 1) { + for (int j = i; j < i + step; j++) { + float a = buf[j], b = buf[j + step]; + buf[j] = a + b; + buf[j + step] = a - b; + } + } + } + for (int i = 0; i < TQ_BLOCK_SIZE; i++) buf[i] *= TQ_INV_SQRT32 * TQ3_0_SIGNS[i]; +} + +/* Nearest centroid for TQ3 (8 centroids) */ +static int tq3_0_choose_index(float val) { + /* Binary search on midpoints of TQ3_0_CENTROIDS */ + if (val < -1.644041f) return 0; + if (val < -1.015870f) return 1; + if (val < -0.493925f) return 2; + if (val < -0.008701f) return 3; + if (val < 0.477664f) return 4; + if (val < 1.001363f) return 5; + if (val < 1.633223f) return 6; + return 7; +} + +/* Nearest centroid for TQ4 (16 centroids) */ +static int tq4_0_choose_index(float val) { + /* Binary search on midpoints of TQ4_0_CENTROIDS */ + if (val < -2.400804f) return 0; + if (val < -1.843532f) return 1; + if (val < -1.437139f) return 2; + if (val < -1.099286f) return 3; + if (val < -0.799550f) return 4; + if (val < -0.522404f) return 5; + if (val < -0.258222f) return 6; + if (val < 0.000000f) return 7; + if (val < 0.258222f) return 8; + if (val < 0.522404f) return 9; + if (val < 0.799550f) return 10; + if (val < 1.099286f) return 11; + if (val < 1.437139f) return 12; + if (val < 1.843532f) return 13; + if (val < 2.400804f) return 14; + return 15; +} + +/* ---------- TQ3_1S quantization ---------- */ + +void quantize_row_tq3_1s_ref(const float * GGML_RESTRICT x, block_tq3_1s * GGML_RESTRICT y, int64_t k) { + assert(k % QK_TQ3_0 == 0); + const int nb = k / QK_TQ3_0; + + for (int block = 0; block < nb; block++) { + const float * src_blk = x + block * QK_TQ3_0; + block_tq3_1s * blk = &y[block]; + + /* 1. Forward RHT */ + float buf[TQ_BLOCK_SIZE]; + memcpy(buf, src_blk, TQ_BLOCK_SIZE * sizeof(float)); + tq3_0_rht_forward(buf); + + /* 2. Split into two halves, compute RMS per half */ + float rms0 = 0.0f, rms1 = 0.0f; + for (int j = 0; j < 16; j++) rms0 += buf[j] * buf[j]; + for (int j = 16; j < 32; j++) rms1 += buf[j] * buf[j]; + rms0 = sqrtf(rms0 / 16.0f); + rms1 = sqrtf(rms1 / 16.0f); + + /* 3. Scale search (9 points) */ + static const float scales[] = { 0.6f, 0.7f, 0.8f, 0.9f, 1.0f, 1.1f, 1.2f, 1.35f, 1.5f }; + float best_d0 = rms0, best_d1 = rms1; + float best_err = 1e30f; + + for (int si = 0; si < 9; si++) { + float d0 = rms0 * scales[si]; + float d1 = rms1 * scales[si]; + float inv0 = (d0 > 1e-10f) ? 1.0f / d0 : 0.0f; + float inv1 = (d1 > 1e-10f) ? 1.0f / d1 : 0.0f; + + float err = 0.0f; + for (int j = 0; j < 16; j++) { + int idx = tq3_0_choose_index(buf[j] * inv0); + float diff = buf[j] - TQ3_0_CENTROIDS[idx] * d0; + err += diff * diff; + } + for (int j = 16; j < 32; j++) { + int idx = tq3_0_choose_index(buf[j] * inv1); + float diff = buf[j] - TQ3_0_CENTROIDS[idx] * d1; + err += diff * diff; + } + if (err < best_err) { + best_err = err; + best_d0 = d0; + best_d1 = d1; + } + } + + /* 4. Iterative refinement (6 iterations) */ + for (int iter = 0; iter < 6; iter++) { + float inv0 = (best_d0 > 1e-10f) ? 1.0f / best_d0 : 0.0f; + float inv1 = (best_d1 > 1e-10f) ? 1.0f / best_d1 : 0.0f; + + float num0 = 0.0f, den0 = 0.0f; + float num1 = 0.0f, den1 = 0.0f; + for (int j = 0; j < 16; j++) { + int idx = tq3_0_choose_index(buf[j] * inv0); + float c = TQ3_0_CENTROIDS[idx]; + num0 += buf[j] * c; + den0 += c * c; + } + for (int j = 16; j < 32; j++) { + int idx = tq3_0_choose_index(buf[j] * inv1); + float c = TQ3_0_CENTROIDS[idx]; + num1 += buf[j] * c; + den1 += c * c; + } + if (den0 > 1e-10f) best_d0 = num0 / den0; + if (den1 > 1e-10f) best_d1 = num1 / den1; + } + + /* 5. Final quantize + pack */ + float inv0 = (best_d0 > 1e-10f) ? 1.0f / best_d0 : 0.0f; + float inv1 = (best_d1 > 1e-10f) ? 1.0f / best_d1 : 0.0f; + + blk->d0 = GGML_FP32_TO_FP16(best_d0); + blk->d1 = GGML_FP32_TO_FP16(best_d1); + memset(blk->qs, 0, QK_TQ3_0 * 3 / 8); + + /* TQ3 packing: 4 groups of 8 indices packed into 3 bytes each */ + for (int g = 0; g < 4; g++) { + uint8_t indices[8]; + for (int i = 0; i < 8; i++) { + int j = g * 8 + i; + float inv = (j < 16) ? inv0 : inv1; + indices[i] = (uint8_t)tq3_0_choose_index(buf[j] * inv); + } + uint8_t * qp = blk->qs + g * 3; + qp[0] = (indices[0] & 7) | ((indices[1] & 7) << 3) | ((indices[2] & 3) << 6); + qp[1] = ((indices[2] >> 2) & 1) | ((indices[3] & 7) << 1) | ((indices[4] & 7) << 4) | ((indices[5] & 1) << 7); + qp[2] = ((indices[5] >> 1) & 3) | ((indices[6] & 7) << 2) | ((indices[7] & 7) << 5); + } + } +} + +void dequantize_row_tq3_1s(const block_tq3_1s * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { + assert(k % QK_TQ3_0 == 0); + const int nb = k / QK_TQ3_0; + + for (int blk_i = 0; blk_i < nb; blk_i++) { + float d0 = GGML_FP16_TO_FP32(x[blk_i].d0); + float d1 = GGML_FP16_TO_FP32(x[blk_i].d1); + + /* Unpack 3-bit indices */ + float buf[32]; + for (int g = 0; g < 4; g++) { + const uint8_t * qp = x[blk_i].qs + g * 3; + uint8_t idx[8]; + idx[0] = qp[0] & 7; + idx[1] = (qp[0] >> 3) & 7; + idx[2] = ((qp[0] >> 6) | (qp[1] << 2)) & 7; + idx[3] = (qp[1] >> 1) & 7; + idx[4] = (qp[1] >> 4) & 7; + idx[5] = ((qp[1] >> 7) | (qp[2] << 1)) & 7; + idx[6] = (qp[2] >> 2) & 7; + idx[7] = (qp[2] >> 5) & 7; + + for (int i = 0; i < 8; i++) { + int j = g * 8 + i; + float d = (j < 16) ? d0 : d1; + buf[j] = TQ3_0_CENTROIDS[idx[i]] * d; + } + } + + /* Inverse RHT */ + tq3_0_rht_inverse(buf); + + memcpy(y + blk_i * QK_TQ3_0, buf, QK_TQ3_0 * sizeof(float)); + } +} + +size_t quantize_tq3_1s(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_TQ3_0 == 0); + + size_t row_size = (n_per_row / QK_TQ3_0) * sizeof(block_tq3_1s); + for (int64_t row = 0; row < nrows; row++) { + quantize_row_tq3_1s_ref( + src + row * n_per_row, + (block_tq3_1s *)((char *)dst + row * row_size), + n_per_row + ); + } + return nrows * row_size; +} + +/* ---------- TQ4_1S quantization ---------- */ + +void quantize_row_tq4_1s_ref(const float * GGML_RESTRICT x, block_tq4_1s * GGML_RESTRICT y, int64_t k) { + assert(k % QK_TQ4_1S == 0); + const int nb = k / QK_TQ4_1S; + + for (int block = 0; block < nb; block++) { + const float * src_blk = x + block * QK_TQ4_1S; + block_tq4_1s * blk = &y[block]; + + /* 1. Forward RHT */ + float buf[TQ_BLOCK_SIZE]; + memcpy(buf, src_blk, TQ_BLOCK_SIZE * sizeof(float)); + tq3_0_rht_forward(buf); + + /* 2. Split into two halves, compute RMS per half */ + float rms0 = 0.0f, rms1 = 0.0f; + for (int j = 0; j < 16; j++) rms0 += buf[j] * buf[j]; + for (int j = 16; j < 32; j++) rms1 += buf[j] * buf[j]; + rms0 = sqrtf(rms0 / 16.0f); + rms1 = sqrtf(rms1 / 16.0f); + + /* 3. Scale search (9 points) */ + static const float scales[] = { 0.6f, 0.7f, 0.8f, 0.9f, 1.0f, 1.1f, 1.2f, 1.35f, 1.5f }; + float best_d0 = rms0, best_d1 = rms1; + float best_err = 1e30f; + + for (int si = 0; si < 9; si++) { + float d0 = rms0 * scales[si]; + float d1 = rms1 * scales[si]; + float inv0 = (d0 > 1e-10f) ? 1.0f / d0 : 0.0f; + float inv1 = (d1 > 1e-10f) ? 1.0f / d1 : 0.0f; + + float err = 0.0f; + for (int j = 0; j < 16; j++) { + int idx = tq4_0_choose_index(buf[j] * inv0); + float diff = buf[j] - TQ4_0_CENTROIDS[idx] * d0; + err += diff * diff; + } + for (int j = 16; j < 32; j++) { + int idx = tq4_0_choose_index(buf[j] * inv1); + float diff = buf[j] - TQ4_0_CENTROIDS[idx] * d1; + err += diff * diff; + } + if (err < best_err) { + best_err = err; + best_d0 = d0; + best_d1 = d1; + } + } + + /* 4. Iterative refinement (6 iterations) */ + for (int iter = 0; iter < 6; iter++) { + float inv0 = (best_d0 > 1e-10f) ? 1.0f / best_d0 : 0.0f; + float inv1 = (best_d1 > 1e-10f) ? 1.0f / best_d1 : 0.0f; + + float num0 = 0.0f, den0 = 0.0f; + float num1 = 0.0f, den1 = 0.0f; + for (int j = 0; j < 16; j++) { + int idx = tq4_0_choose_index(buf[j] * inv0); + float c = TQ4_0_CENTROIDS[idx]; + num0 += buf[j] * c; + den0 += c * c; + } + for (int j = 16; j < 32; j++) { + int idx = tq4_0_choose_index(buf[j] * inv1); + float c = TQ4_0_CENTROIDS[idx]; + num1 += buf[j] * c; + den1 += c * c; + } + if (den0 > 1e-10f) best_d0 = num0 / den0; + if (den1 > 1e-10f) best_d1 = num1 / den1; + } + + /* 5. Final quantize + pack (nibble packing) */ + float inv0 = (best_d0 > 1e-10f) ? 1.0f / best_d0 : 0.0f; + float inv1 = (best_d1 > 1e-10f) ? 1.0f / best_d1 : 0.0f; + + blk->d0 = GGML_FP32_TO_FP16(best_d0); + blk->d1 = GGML_FP32_TO_FP16(best_d1); + memset(blk->qs, 0, QK_TQ4_1S / 2); + + for (int j = 0; j < QK_TQ4_1S; j++) { + float inv = (j < 16) ? inv0 : inv1; + int idx = tq4_0_choose_index(buf[j] * inv); + blk->qs[j / 2] |= (uint8_t)((idx & 0xF) << ((j & 1) * 4)); + } + } +} + +void dequantize_row_tq4_1s(const block_tq4_1s * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { + assert(k % QK_TQ4_1S == 0); + const int nb = k / QK_TQ4_1S; + + for (int blk_i = 0; blk_i < nb; blk_i++) { + float d0 = GGML_FP16_TO_FP32(x[blk_i].d0); + float d1 = GGML_FP16_TO_FP32(x[blk_i].d1); + + float buf[32]; + for (int j = 0; j < 32; j++) { + uint8_t idx = (x[blk_i].qs[j / 2] >> ((j & 1) * 4)) & 0xF; + float d = (j < 16) ? d0 : d1; + buf[j] = TQ4_0_CENTROIDS[idx] * d; + } + + /* Inverse RHT */ + tq3_0_rht_inverse(buf); + + memcpy(y + blk_i * QK_TQ4_1S, buf, QK_TQ4_1S * sizeof(float)); + } +} + +size_t quantize_tq4_1s(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_TQ4_1S == 0); + + size_t row_size = (n_per_row / QK_TQ4_1S) * sizeof(block_tq4_1s); + for (int64_t row = 0; row < nrows; row++) { + quantize_row_tq4_1s_ref( + src + row * n_per_row, + (block_tq4_1s *)((char *)dst + row * row_size), + n_per_row + ); + } + return nrows * row_size; +} diff --git a/llama.cpp/ggml/src/ggml.c b/llama.cpp/ggml/src/ggml.c --- a/llama.cpp/ggml/src/ggml.c +++ b/llama.cpp/ggml/src/ggml.c @@ -674,6 +674,55 @@ static const struct ggml_type_traits type_traits[GGML_TYPE_COUNT] = { .to_float = (ggml_to_float_t) dequantize_row_q1_0, .from_float_ref = (ggml_from_float_t) quantize_row_q1_0_ref, }, + // opencoti-hook: turboquant — see docs/features/poly_kv.md (M6-S2) + [GGML_TYPE_TURBO3_0] = { + .type_name = "turbo3", + .blck_size = QK_TURBO3, + .type_size = sizeof(block_turbo3_0), + .is_quantized = true, + .to_float = (ggml_to_float_t) dequantize_row_turbo3_0, + .from_float_ref = (ggml_from_float_t) quantize_row_turbo3_0_ref, + }, + [GGML_TYPE_TURBO4_0] = { + .type_name = "turbo4", + .blck_size = QK_TURBO4, + .type_size = sizeof(block_turbo4_0), + .is_quantized = true, + .to_float = (ggml_to_float_t) dequantize_row_turbo4_0, + .from_float_ref = (ggml_from_float_t) quantize_row_turbo4_0_ref, + }, + [GGML_TYPE_TURBO2_0] = { + .type_name = "turbo2", + .blck_size = QK_TURBO2, + .type_size = sizeof(block_turbo2_0), + .is_quantized = true, + .to_float = (ggml_to_float_t) dequantize_row_turbo2_0, + .from_float_ref = (ggml_from_float_t) quantize_row_turbo2_0_ref, + }, + [GGML_TYPE_TURBO8_0] = { + .type_name = "turbo8", + .blck_size = QK_TURBO8, + .type_size = sizeof(block_turbo8_0), + .is_quantized = true, + .to_float = (ggml_to_float_t) dequantize_row_turbo8_0, + .from_float_ref = (ggml_from_float_t) quantize_row_turbo8_0_ref, + }, + [GGML_TYPE_TQ3_1S] = { + .type_name = "tq3_1s", + .blck_size = QK_TQ3_0, + .type_size = sizeof(block_tq3_1s), + .is_quantized = true, + .to_float = (ggml_to_float_t) dequantize_row_tq3_1s, + .from_float_ref = (ggml_from_float_t) quantize_row_tq3_1s_ref, + }, + [GGML_TYPE_TQ4_1S] = { + .type_name = "tq4_1s", + .blck_size = QK_TQ4_1S, + .type_size = sizeof(block_tq4_1s), + .is_quantized = true, + .to_float = (ggml_to_float_t) dequantize_row_tq4_1s, + .from_float_ref = (ggml_from_float_t) quantize_row_tq4_1s_ref, + }, [GGML_TYPE_Q4_0] = { .type_name = "q4_0", .blck_size = QK4_0, @@ -1081,9 +1130,10 @@ static const char * GGML_OP_NAME[GGML_OP_COUNT] = { "MOE_FUSED_UP_GATE", "STREAMING_FLASH_ATTN", "FLASH_ATTN_TAIL_PARTIAL", + "TURBO_WHT", }; -static_assert(GGML_OP_COUNT == 99, "GGML_OP_COUNT != 99"); +static_assert(GGML_OP_COUNT == 100, "GGML_OP_COUNT != 100"); static const char * GGML_OP_SYMBOL[GGML_OP_COUNT] = { "none", @@ -1194,9 +1244,10 @@ static const char * GGML_OP_SYMBOL[GGML_OP_COUNT] = { "moe_up_gate(up,gate,x)", "streaming_flash_attn(q,k,v)", "flash_attn_tail_partial(q,k,v)", + "turbo_wht(x)", }; -static_assert(GGML_OP_COUNT == 99, "GGML_OP_COUNT != 99"); +static_assert(GGML_OP_COUNT == 100, "GGML_OP_COUNT != 100"); static_assert(GGML_OP_POOL_COUNT == 2, "GGML_OP_POOL_COUNT != 2"); @@ -5577,6 +5628,40 @@ struct ggml_tensor * ggml_flash_attn_ext_tail_partial( return result; } +// opencoti-hook: turboquant perf-lift (Level B) — F5 M6-S2 #391. Walsh–Hadamard +// rotation op for the fused turbo FA-VEC path (Q forward-rotate / FA-output +// inverse-rotate). B-adapted vs the AtomicBot fork: `scale` is normally NULL and +// the CUDA kernel resolves InnerQ from the device symbol (set-rows.cu), so no +// memory-context graph-tensor plumbing is needed; a non-NULL `scale` (fork-compat) +// is honoured as src[1] when supplied. Output shares a's shape/type (f32). +struct ggml_tensor * ggml_turbo_wht( + struct ggml_context * ctx, + struct ggml_tensor * a, + int direction, + int group_size, + struct ggml_tensor * scale) { + GGML_ASSERT(ggml_is_contiguous(a)); + GGML_ASSERT(a->type == GGML_TYPE_F32); + GGML_ASSERT(direction == 0 || direction == 1); + + if (group_size == 0) { + group_size = (a->ne[0] % 128 == 0) ? 128 : 64; + } + GGML_ASSERT(group_size == 32 || group_size == 64 || group_size == 128); + GGML_ASSERT(a->ne[0] % group_size == 0); + + struct ggml_tensor * result = ggml_new_tensor(ctx, GGML_TYPE_F32, 4, a->ne); + + result->op = GGML_OP_TURBO_WHT; + result->src[0] = a; + result->src[1] = scale; // InnerQ scale_inv (NULL ⇒ CUDA resolves from device symbol) + + memcpy((char *) result->op_params + 0, &direction, sizeof(int)); + memcpy((char *) result->op_params + sizeof(int), &group_size, sizeof(int)); + + return result; +} + void ggml_flash_attn_ext_set_prec( struct ggml_tensor * a, enum ggml_prec prec) { @@ -7923,6 +8008,13 @@ size_t ggml_quantize_chunk( case GGML_TYPE_IQ1_M: result = quantize_iq1_m (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_IQ4_NL: result = quantize_iq4_nl (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_IQ4_XS: result = quantize_iq4_xs (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; + // opencoti-hook: turboquant — see docs/features/poly_kv.md (M6-S2) + case GGML_TYPE_TURBO3_0: result = quantize_turbo3_0(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; + 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; + 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: { size_t elemsize = sizeof(ggml_fp16_t); diff --git a/llama.cpp/src/llama-graph.cpp b/llama.cpp/src/llama-graph.cpp --- a/llama.cpp/src/llama-graph.cpp +++ b/llama.cpp/src/llama-graph.cpp @@ -1988,6 +1988,57 @@ ggml_tensor * llm_graph_context::build_attn_mha( ggml_tensor * v_mla, float kq_scale, int il) const { + // opencoti-hook: turboquant perf-lift (M6-S2 Level B) — see docs/features/poly_kv.md + + // plan lexical-gliding-flurry. FUSED FA-VEC turbo path: for turbo2/3/4 with K==V at + // head_dim ∈ {128,256} (the vec-instanced widths) we read the COMPRESSED cache in-register + // inside flash-attn instead of materializing it to f16. The expensive Walsh–Hadamard rotation + // then runs only on the n_q-sized Q (forward) and the small FA OUTPUT (inverse) — never on the + // KV cache. InnerQ ×scale_inv is folded into BOTH ggml_turbo_wht kernels via the device-resident + // scale symbol keyed on ne[0]=head_dim (B-adapted; no graph-tensor scale src). Parseval identity: + // WHT(Q⊘s)·K_stored == Q_true·K_true and WHT⁻¹(Σp·V_stored)⊘s == Σp·V_true, so attention is exact + // in true space (logit-equiv gate, NOT byte-equality). + // + // FALLBACK = Level-A materialize (the else-branch below): turbo8 (no vec tier), head_dim 512 + // global layers (no D=512 turbo vec instance), asymmetric K≠V, or the non-flash-attn branch — + // 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; + }; + // 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 + // materialize) but LOSES at PREFILL (large n_q rides the vec kernel, not MMA: ~817-1537 vs ~2500 + // tps). So fuse only for decode-like ubatches (small per-stream query count); large prefill + // ubatches fall through to the Level-A materialize→f16→MMA path. Threshold is the per-stream + // query-token count = q->ne[2] / n_stream (decode=1, prefill=chunk). Env-tunable for the A/B: + // TBV_FUSE_MAX_NQ = max per-stream n_q that still fuses (default 8). Set huge → always-fuse + // (= L4 baseline); set 0 → always-materialize (= Level-A). No rebuild needed to sweep. + static const int tbv_fuse_max_nq = []() { + const char * e = getenv("TBV_FUSE_MAX_NQ"); + return e ? atoi(e) : 8; + }(); + const int64_t tbv_n_stream = k->ne[3] > 0 ? k->ne[3] : 1; + const int64_t tbv_nq_per_strm = q->ne[2] / tbv_n_stream; + const bool fused_turbo = + cparams.flash_attn && kq_b == nullptr && + is_turbo_vec(k->type) && k->type == v->type && + (k->ne[0] == 128 || k->ne[0] == 256) && + tbv_nq_per_strm <= tbv_fuse_max_nq; + + if (fused_turbo) { + // forward-rotate Q (×scale_inv → signs₁ → WHT → 1/√G·signs₂); group auto-picks 128 since + // head_dim%128==0 for both 128 and 256. Q reaches here as the contiguous projection reshape, + // but guard defensively (a rope'd view would trip the op's contiguity assert). + if (!ggml_is_contiguous(q)) { q = ggml_cont(ctx0, q); } + q = ggml_turbo_wht(ctx0, q, /*direction=*/0, /*group_size=*/0, /*scale=*/nullptr); + } else { + const auto is_turbo = [](enum ggml_type t) { + return t == GGML_TYPE_TURBO2_0 || t == GGML_TYPE_TURBO3_0 || t == GGML_TYPE_TURBO4_0 || t == GGML_TYPE_TURBO8_0; + }; + if (is_turbo(k->type)) { k = ggml_cast(ctx0, k, GGML_TYPE_F16); } + if (is_turbo(v->type)) { v = ggml_cast(ctx0, v, GGML_TYPE_F16); } + } + const bool v_trans = v->nb[1] > v->nb[2]; // split the batch into streams if needed @@ -2025,6 +2076,14 @@ ggml_tensor * llm_graph_context::build_attn_mha( ggml_flash_attn_ext_add_sinks(cur, sinks); ggml_flash_attn_ext_set_prec (cur, GGML_PREC_F32); + if (fused_turbo) { + // opencoti-hook: turboquant perf-lift (M6-S2 Level B). Inverse-WHT + InnerQ un-equalize + // on the SMALL FA output (signs₂ → WHT → 1/√G·signs₁ → ×scale_inv), recovering true-space + // attention from the rotated/equalized cache. ne[0]=head_dim_v keys the InnerQ slot; for + // turbo K==V the per-width scale matches the one folded into the forward-Q rotation above. + cur = ggml_turbo_wht(ctx0, cur, /*direction=*/1, /*group_size=*/0, /*scale=*/nullptr); + } + if (v_mla) { #if 0 // v_mla can be applied as a matrix-vector multiplication with broadcasting across dimension 3 == n_tokens. diff --git a/llamafile/build-functions.sh b/llamafile/build-functions.sh --- a/llamafile/build-functions.sh +++ b/llamafile/build-functions.sh @@ -143,6 +143,22 @@ done fi + # 3b. opencoti-hook: turboquant perf-lift (Level B) — F5 M6-S2 #392; see 0073-turboquant-kv + # in vendors/patches/llamafile/ + docs/protocols/UPSTREAM_SYNC.md. + # TurboQuant KV uses a FUSED in-register FA-VEC read (no dequant-on-lift), so the turbo + # K==V vec instances must compile in the DEFAULT build too — not just under FA_ALL_QUANTS. + # Omitting them leaves the DSO without ggml_cuda_flash_attn_ext_vec_case; + # dlopen then fails the instant turbo KV reaches flash-attention (bug: FA-VEC build-list, + # sibling of bug-339). Guarded so it is a no-op when FA_ALL_QUANTS=1 (the glob above + # already globbed these) and on a vanilla tree (the per-file [ -f ] test). + 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 + [ -f "$f" ] && CUDA_SOURCES="$CUDA_SOURCES $f" + done + fi + # 4. mmf instances (always included) for f in "$ti_dir"/mmf-*.cu; do [ -f "$f" ] && CUDA_SOURCES="$CUDA_SOURCES $f" @@ -228,6 +244,7 @@ $llama_cpp_dir/ggml/src/ggml-neo-pipeline.cpp \ $llama_cpp_dir/ggml/src/ggml-backend-meta.cpp \ $llama_cpp_dir/ggml/src/ggml-quants.c \ + $llama_cpp_dir/ggml/src/ggml-turbo-quant.c \ $llama_cpp_dir/ggml/src/ggml-threading.cpp" echo "Compiling core GGML sources..."