From e65cf46408fa52513fcaf6a522349c58ad937022 Mon Sep 17 00:00:00 2001 From: NeuralNotwerk Date: Thu, 16 Jul 2026 06:06:13 +0000 Subject: [PATCH] =?UTF-8?q?KV8=3D1=20=E2=80=94=20fp8=20e4m3=20latent=20KV?= =?UTF-8?q?=20cache=20(CPU=20+=20CUDA)=20with=20.coli=5Fkv=20v2?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Store the MLA latent rows (Lc kv_lora + Rc qk_rope) as fp8 e4m3 bytes with a per-row f32 amax/448 scale instead of f32: ~3.9x less KV RAM (a 256k slot drops 47.7 -> ~13 GB), and kv_pool_bytes reports the smaller pool so cap_for_ram/expert_avail stop demoting experts to disk under multi-slot KV. Same regime as DeepSeek-V3's fp8 latent-KV practice: the latent is both key and value in the absorbed attention; per-row scaling keeps e4m3 safe. - c/kv_fp8.h: 256-entry decode LUT + RNE bit-math encoder (saturating at +/-448, signed zero preserved, NaN stored as inert 0), per-row quantizer with a subnormal-amax guard (448/amax would overflow to +inf). - attention_rows: producer quantizes the fresh normed/roped row in place of the f32 memcpy; absorb score/context loops LUT-dequant inline with the scale hoisted out of the dot; the non-absorb prefill dequants Lc into an f32 staging row for the kv_b matmul and reads Rc via the LUT. - CUDA: attention_absorb_kernel8/_batch_kernel8 (hw fp8 cvt, sm_89+) behind coli_cuda_attention_absorb8/_absorb_batch8/_project_batch8 — 1/4 the PCIe traffic of the f32 uploads; wired at all three COLI_CUDA_ATTN call sites and into the Windows DLL loader. COLI_CUDA_PIPE and Metal still read f32 rows, so KV8 forces the pipe off and auto-disables under Metal. - .coli_kv v2 (COLIKV2 magic, dtype in the header) through the persistent- handle + staging-buffer append path: fp8 rows + scales at ~46 KB/token; v1 files quantize on resume and stay intact on disk until the first append rewrites them as v2 (magic self-heal in kv_disk_open), so a crash before the first save loses nothing. - Tests: exhaustive e4m3 roundtrip + RNE ties incl. the binade-boundary mantissa carry, KV8 kv_alloc transitions, v1/v2 disk round-trip/upgrade/ reject/self-heal, CUDA fp8 kernel checks vs dequantized-host reference. Validation: tiny-oracle KV8=0 bit-identical (TF 32/32, gen 20/20, CPU and CUDA); KV8=1 flips only two positions that transformers' own f32 forward flips across versions (near-tie margins on a degenerate random-weight model); CPU and CUDA fp8 paths agree token-for-token; ASan/UBSan clean. Real-model A/B (GLM-5.2 744B int4, 3x ~450-token log-lik requests): ppl deltas +2.4% / +0.6% / -3.7% — net ~0, inside the int4 noise floor. Co-Authored-By: Claude Fable 5 --- .gitignore | 6 + README.md | 1 + c/Makefile | 12 +- c/backend_cuda.cu | 146 ++++++++++++++++++++ c/backend_cuda.h | 18 +++ c/backend_loader.c | 24 ++++ c/decode_batch.h | 7 + c/glm.c | 253 +++++++++++++++++++++++++++++------ c/kv_fp8.h | 65 +++++++++ c/tests/test_backend_cuda.cu | 40 ++++++ c/tests/test_kv_alloc.c | 22 ++- c/tests/test_kv_disk.c | 102 ++++++++++++++ c/tests/test_kv_fp8.c | 88 ++++++++++++ docs/ENVIRONMENT.md | 1 + 14 files changed, 737 insertions(+), 48 deletions(-) create mode 100644 c/kv_fp8.h create mode 100644 c/tests/test_kv_disk.c create mode 100644 c/tests/test_kv_fp8.c diff --git a/.gitignore b/.gitignore index 981e74005..c02c134ed 100644 --- a/.gitignore +++ b/.gitignore @@ -67,4 +67,10 @@ c/bench/ c/tests/test_decode_batch c/tests/test_i4_acc512 c/tests/test_idot +c/tests/test_kv_alloc +c/tests/test_kv_alloc.exe +c/tests/test_kv_fp8 +c/tests/test_kv_fp8.exe +c/tests/test_kv_disk +c/tests/test_kv_disk.exe c/tests/test_uring diff --git a/README.md b/README.md index 2e2204805..40583f7e3 100644 --- a/README.md +++ b/README.md @@ -66,6 +66,7 @@ The engine is a single C file (`c/glm.c`) plus small headers. No BLAS, no Python - **Quantization kernels**: int8 / packed int4 / packed int2, per-row scales, AVX2, dequant-on-use. Packing validated bit-identical to the int8 container. - **DSA sparse attention** — GLM-5.2's lightning indexer, faithful to the reference `glm_moe_dsa` modeling: per-layer top-2048 causal key selection (full/shared indexer layers), auto-detected from the `out-idx-*` weights (`--indexer` converter mode, ~189 MB extracted from the FP8 repo). Validated exact: forcing the selection to keep every key reproduces dense attention token-for-token. `DSA=0` disables, `DSA_TOPK` overrides. - **KV-cache persistence** — conversations reopen **warm** across engine restarts: serve mode appends the compressed MLA KV to `.coli_kv` after every turn (~182 KB/token, crash-safe) and resumes it at startup with zero re-prefill. Validated byte-identical to an uninterrupted session. `KVSAVE=0` disables. +- **FP8 KV cache** (`KV8=1`, CPU attention path) — the MLA latent rows are stored as fp8 e4m3 with a per-row scale: ~3.9× less KV RAM (a 256k-token context drops from ~47.7 to ~13 GB), `.coli_kv` shrinks to ~46 KB/token (v2 format; f32 saves are quantized on resume and rewritten), and the RAM budget planner sees the smaller pool, so multi-slot serving stops demoting experts to disk. DeepSeek-V3-class KV quantization noise; covered on the CPU and CUDA attention paths (`COLI_CUDA_ATTN=1` uses fp8 absorb kernels at ¼ the PCIe traffic); auto-disabled under `COLI_METAL`. - **Router-lookahead prefetch** (`PILOT=1`, experimental) — the next layer's routing is 71.6% predictable from the current layer's post-attention state (measured); a dedicated I/O thread prefetches those experts while the current layer computes. - **Batch-union MoE**: in prefill (and MTP verification), each unique expert of the batch is read once and applied to every position that routes to it. - **Byte-level BPE tokenizer in C** (GPT-2-style with Unicode-property regex, 320k merges). diff --git a/c/Makefile b/c/Makefile index 1ed6df6b4..6f2a167ef 100644 --- a/c/Makefile +++ b/c/Makefile @@ -166,7 +166,7 @@ else PYTHON ?= python3 endif CUDA_OBJ = -TEST_BINS = tests/test_json$(EXE) tests/test_st$(EXE) tests/test_tier$(EXE) tests/test_grammar$(EXE) tests/test_schema_gbnf$(EXE) tests/test_decode_batch$(EXE) tests/test_idot$(EXE) tests/test_kv_alloc$(EXE) tests/test_i4_acc512$(EXE) tests/test_compat_direct$(EXE) +TEST_BINS = tests/test_json$(EXE) tests/test_st$(EXE) tests/test_tier$(EXE) tests/test_grammar$(EXE) tests/test_schema_gbnf$(EXE) tests/test_decode_batch$(EXE) tests/test_idot$(EXE) tests/test_kv_alloc$(EXE) tests/test_kv_fp8$(EXE) tests/test_kv_disk$(EXE) tests/test_i4_acc512$(EXE) tests/test_compat_direct$(EXE) ifneq (,$(LINUX)) TEST_BINS += tests/test_uring$(EXE) endif @@ -212,7 +212,7 @@ all: glm$(EXE) # phony 'glm' → 'glm.exe' on Windows (so 'make glm' and 'coli build' work on every platform) glm: glm$(EXE) -glm$(EXE): glm.c st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h $(CUDA_OBJ) $(METAL_OBJ) +glm$(EXE): glm.c st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h decode_batch.h kv_fp8.h $(CUDA_OBJ) $(METAL_OBJ) $(CC) $(CFLAGS) glm.c $(CUDA_OBJ) $(METAL_OBJ) -o glm$(EXE) $(LDFLAGS) # Windows runtime loader object: resolves coli_cuda_* from coli_cuda.dll. @@ -297,7 +297,13 @@ tests/test_decode_batch$(EXE): tests/test_decode_batch.c decode_batch.h tests/test_idot$(EXE): tests/test_idot.c glm.c st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) -tests/test_kv_alloc$(EXE): tests/test_kv_alloc.c glm.c st.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h +tests/test_kv_alloc$(EXE): tests/test_kv_alloc.c glm.c st.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h kv_fp8.h + $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) + +tests/test_kv_fp8$(EXE): tests/test_kv_fp8.c kv_fp8.h decode_batch.h + $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) + +tests/test_kv_disk$(EXE): tests/test_kv_disk.c glm.c st.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h kv_fp8.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) tests/test_i4_acc512$(EXE): tests/test_i4_acc512.c diff --git a/c/backend_cuda.cu b/c/backend_cuda.cu index 9ce142fbf..371d8a2cc 100644 --- a/c/backend_cuda.cu +++ b/c/backend_cuda.cu @@ -1,6 +1,7 @@ #include "backend_cuda.h" #include +#include /* KV8: fp8 e4m3 latent KV (hw cvt on sm_89+, our arches) */ #include #include @@ -25,6 +26,7 @@ typedef struct { size_t qx_cap, qscale_cap; float *host_x,*host_y; size_t host_x_cap,host_y_cap; float *aq,*al,*ar,*ac; size_t aq_cap,al_cap,ar_cap,ac_cap; + float *alsc,*arsc; size_t alsc_cap,arsc_cap; /* KV8: scale per-riga di latent/rope */ float *pipe_buf[24]; size_t pipe_cap[24]; /* scratch persistenti del resident pipeline */ cudaStream_t stream; void *group_desc; size_t group_desc_cap; @@ -287,6 +289,12 @@ __global__ static void grouped_down_w4(float *y,const float *x,const GroupDesc * if(!threadIdx.x)y[(size_t)(d.offset+s)*D+o]=p[0]*d.ds[o]; } +/* KV8: decodifica un byte e4m3 (cvt hardware da sm_89; PTX/software prima). La scala + * per-riga viaggia in un array f32 separato e si applica una volta per score/peso. */ +__device__ static inline float fp8_e4m3(uint8_t b){ + return __half2float(__half(__nv_cvt_fp8_to_halfraw((__nv_fp8_storage_t)b,__NV_E4M3))); +} + __global__ static void attention_absorb_kernel(float *ctx,const float *q,const float *latent, const float *rope,const void *weights,const float *wscale, int fmt,int H,int Q,int R,int V,int K,int T,float scale){ @@ -338,6 +346,68 @@ __global__ static void attention_absorb_batch_kernel(float *ctx,const float *q, ctx[((size_t)s*H+h)*V+v]=a*(fmt?wscale[row]:1.f);} } +/* ---- KV8: gemelli fp8 dei due kernel di assorbimento. latent/rope arrivano come + * byte e4m3 + scala f32 per riga. Stessa matematica dei gemelli f32: la scala esce + * dal dot (score = Lsc·Σ q·v + Rsc·Σ qr·v) e per il contesto si fonde nel peso + * softmax, cosi' i loop interni restano una FMA per byte. */ +__global__ static void attention_absorb_kernel8(float *ctx,const float *q,const uint8_t *latent, + const float *lsc,const uint8_t *rope,const float *rsc,const void *weights, + const float *wscale,int fmt,int H,int Q,int R,int V,int K,int T,float scale){ + int h=blockIdx.x,tid=threadIdx.x,rbase=h*(Q+V);extern __shared__ float sm[]; + float *qa=sm,*cl=qa+K,*scores=cl+K; + for(int k=tid;k=S||nt<1)return; + extern __shared__ float sm[];float *qa=sm,*cl=qa+K,*scores=cl+K,*red=scores+T; + const float *qs=q+((size_t)s*H+h)*(Q+R); + for(int k=tid;k>1;n;n>>=1){if(tid>1;n;n>>=1){if(tid= bytes) return 1; if (*ptr) cudaFree(*ptr); @@ -403,6 +473,7 @@ extern "C" void coli_cuda_shutdown(void) { if (ctx->qx) cudaFree(ctx->qx); if (ctx->qscale) cudaFree(ctx->qscale); if(ctx->aq)cudaFree(ctx->aq);if(ctx->al)cudaFree(ctx->al);if(ctx->ar)cudaFree(ctx->ar);if(ctx->ac)cudaFree(ctx->ac); + if(ctx->alsc)cudaFree(ctx->alsc);if(ctx->arsc)cudaFree(ctx->arsc); for(int b=0;b<24;b++) if(ctx->pipe_buf[b]) cudaFree(ctx->pipe_buf[b]); if (ctx->host_x) cudaFreeHost(ctx->host_x); if (ctx->host_y) cudaFreeHost(ctx->host_y); @@ -415,6 +486,7 @@ extern "C" void coli_cuda_shutdown(void) { ctx->x_cap = ctx->y_cap = ctx->gate_cap = ctx->up_cap = 0; ctx->qx_cap=ctx->qscale_cap=0; ctx->aq_cap=ctx->al_cap=ctx->ar_cap=ctx->ac_cap=0; + ctx->alsc=ctx->arsc=nullptr; ctx->alsc_cap=ctx->arsc_cap=0; ctx->host_x_cap=ctx->host_y_cap=0; ctx->group_desc=nullptr; ctx->group_desc_cap=0; } @@ -770,6 +842,80 @@ extern "C" int coli_cuda_attention_project_batch(ColiCudaTensor *w,ColiCudaTenso return attention_absorb_batch_run(w,proj,out,q,latent,rope,S,H,Q,R,V,K,T,scale); } +/* ---- KV8: entry point fp8. Stessi contratti dei gemelli f32; latent/rope viaggiano + * come byte + scala per riga (1/4 del traffico PCIe — a T lunghi e' il collo). */ +extern "C" int coli_cuda_attention_absorb8(ColiCudaTensor *w,float *ctx,const float *q, + const uint8_t *latent,const float *lsc,const uint8_t *rope,const float *rsc, + int H,int Q,int R,int V,int K,int T,float scale){ + if(!w||!ctx||!q||!latent||!lsc||!rope||!rsc||H<1||Q<1||R<1||V<1||K<1||K>512||T<1||T>4096|| + w->I!=K||w->O!=H*(Q+V))return 0; + DeviceContext *dc=find_ctx(w->device);if(!select_ctx(dc))return 0; + size_t qb=(size_t)H*(Q+R)*sizeof(float),lb=(size_t)T*K,rb=(size_t)T*R; + size_t sb=(size_t)T*sizeof(float),cb=(size_t)H*V*sizeof(float); + if(!reserve(&dc->aq,&dc->aq_cap,qb)||!reserve(&dc->al,&dc->al_cap,lb)|| + !reserve(&dc->ar,&dc->ar_cap,rb)||!reserve(&dc->ac,&dc->ac_cap,cb)|| + !reserve(&dc->alsc,&dc->alsc_cap,sb)||!reserve(&dc->arsc,&dc->arsc_cap,sb))return 0; + if(!cuda_ok(cudaMemcpyAsync(dc->aq,q,qb,cudaMemcpyHostToDevice,dc->stream),"attention q upload")|| + !cuda_ok(cudaMemcpyAsync(dc->al,latent,lb,cudaMemcpyHostToDevice,dc->stream),"attention fp8 latent upload")|| + !cuda_ok(cudaMemcpyAsync(dc->alsc,lsc,sb,cudaMemcpyHostToDevice,dc->stream),"attention latent scale upload")|| + !cuda_ok(cudaMemcpyAsync(dc->ar,rope,rb,cudaMemcpyHostToDevice,dc->stream),"attention fp8 rope upload")|| + !cuda_ok(cudaMemcpyAsync(dc->arsc,rsc,sb,cudaMemcpyHostToDevice,dc->stream),"attention rope scale upload"))return 0; + size_t shared=(size_t)(2*K+T)*sizeof(float); + attention_absorb_kernel8<<stream>>>(dc->ac,dc->aq,(const uint8_t*)dc->al, + dc->alsc,(const uint8_t*)dc->ar,dc->arsc,w->weights,w->scales,w->fmt,H,Q,R,V,K,T,scale); + if(!cuda_ok(cudaGetLastError(),"attention absorb8 launch")|| + !cuda_ok(cudaMemcpyAsync(ctx,dc->ac,cb,cudaMemcpyDeviceToHost,dc->stream),"attention context download")|| + !cuda_ok(cudaStreamSynchronize(dc->stream),"attention synchronize"))return 0; + return 1; +} + +static int attention_absorb_batch_run8(ColiCudaTensor *w,ColiCudaTensor *proj,float *out, + const float *q,const uint8_t *latent,const float *lsc,const uint8_t *rope, + const float *rsc,int S,int H,int Q,int R,int V,int K,int T,float scale){ + if(!w||!out||!q||!latent||!lsc||!rope||!rsc||S<1||H<1||Q<1||R<1||V<1||K<1||K>512|| + T8192||w->I!=K||w->O!=H*(Q+V))return 0; + if(proj&&(proj->device!=w->device||proj->I!=H*V))return 0; + DeviceContext *dc=find_ctx(w->device);if(!select_ctx(dc))return 0; + size_t qb=(size_t)S*H*(Q+R)*sizeof(float),lb=(size_t)T*K,rb=(size_t)T*R; + size_t sb=(size_t)T*sizeof(float),cb=(size_t)S*H*V*sizeof(float); + if(!reserve(&dc->aq,&dc->aq_cap,qb)||!reserve(&dc->al,&dc->al_cap,lb)|| + !reserve(&dc->ar,&dc->ar_cap,rb)||!reserve(&dc->ac,&dc->ac_cap,cb)|| + !reserve(&dc->alsc,&dc->alsc_cap,sb)||!reserve(&dc->arsc,&dc->arsc_cap,sb))return 0; + if(!cuda_ok(cudaMemcpyAsync(dc->aq,q,qb,cudaMemcpyHostToDevice,dc->stream),"attention batch q upload")|| + !cuda_ok(cudaMemcpyAsync(dc->al,latent,lb,cudaMemcpyHostToDevice,dc->stream),"attention batch fp8 latent upload")|| + !cuda_ok(cudaMemcpyAsync(dc->alsc,lsc,sb,cudaMemcpyHostToDevice,dc->stream),"attention batch latent scale upload")|| + !cuda_ok(cudaMemcpyAsync(dc->ar,rope,rb,cudaMemcpyHostToDevice,dc->stream),"attention batch fp8 rope upload")|| + !cuda_ok(cudaMemcpyAsync(dc->arsc,rsc,sb,cudaMemcpyHostToDevice,dc->stream),"attention batch rope scale upload"))return 0; + size_t shared=(size_t)(2*K+T+256)*sizeof(float); + attention_absorb_batch_kernel8<<stream>>>(dc->ac,dc->aq, + (const uint8_t*)dc->al,dc->alsc,(const uint8_t*)dc->ar,dc->arsc, + w->weights,w->scales,w->fmt,S,H,Q,R,V,K,T,scale); + if(!cuda_ok(cudaGetLastError(),"attention batch8 launch"))return 0; + const float *src=dc->ac;size_t ob=cb; + if(proj){ + ob=(size_t)S*proj->O*sizeof(float);if(!reserve(&dc->y,&dc->y_cap,ob))return 0; + quant_matmul<<O,S),256,0,dc->stream>>>(dc->y,dc->ac,proj->weights, + proj->scales,proj->fmt,S,proj->I,proj->O,row_bytes(proj->fmt,proj->I)); + if(!cuda_ok(cudaGetLastError(),"attention o_proj launch"))return 0;src=dc->y; + } + if(!cuda_ok(cudaMemcpyAsync(out,src,ob,cudaMemcpyDeviceToHost,dc->stream), + proj?"attention projected output download":"attention batch context download")|| + !cuda_ok(cudaStreamSynchronize(dc->stream),"attention batch synchronize"))return 0; + return 1; +} + +extern "C" int coli_cuda_attention_absorb_batch8(ColiCudaTensor *w,float *ctx,const float *q, + const uint8_t *latent,const float *lsc,const uint8_t *rope,const float *rsc, + int S,int H,int Q,int R,int V,int K,int T,float scale){ + return attention_absorb_batch_run8(w,nullptr,ctx,q,latent,lsc,rope,rsc,S,H,Q,R,V,K,T,scale); +} + +extern "C" int coli_cuda_attention_project_batch8(ColiCudaTensor *w,ColiCudaTensor *proj, + float *out,const float *q,const uint8_t *latent,const float *lsc,const uint8_t *rope, + const float *rsc,int S,int H,int Q,int R,int V,int K,int T,float scale){ + return attention_absorb_batch_run8(w,proj,out,q,latent,lsc,rope,rsc,S,H,Q,R,V,K,T,scale); +} + extern "C" void coli_cuda_tensor_free(ColiCudaTensor *tensor) { if (!tensor) return; DeviceContext *ctx = find_ctx(tensor->device); diff --git a/c/backend_cuda.h b/c/backend_cuda.h index acbe4bc8b..afe68f044 100644 --- a/c/backend_cuda.h +++ b/c/backend_cuda.h @@ -92,6 +92,24 @@ COLI_CUDA_DLLEXPORT int coli_cuda_attention_project_batch(ColiCudaTensor *kv_b,C const float *rope,int S,int H,int Q,int R, int V,int K,int T,float attention_scale); +/* KV8 twins: latent/rope as fp8 e4m3 bytes + one f32 amax/448 scale per row + * (per token, Lc and Rc separately) — 1/4 the PCIe traffic of the f32 paths. */ +COLI_CUDA_DLLEXPORT int coli_cuda_attention_absorb8(ColiCudaTensor *kv_b,float *ctx,const float *q, + const uint8_t *latent,const float *lsc, + const uint8_t *rope,const float *rsc,int H,int Q, + int R,int V,int K,int T,float attention_scale); +COLI_CUDA_DLLEXPORT int coli_cuda_attention_absorb_batch8(ColiCudaTensor *kv_b,float *ctx,const float *q, + const uint8_t *latent,const float *lsc, + const uint8_t *rope,const float *rsc,int S, + int H,int Q,int R,int V,int K,int T, + float attention_scale); +COLI_CUDA_DLLEXPORT int coli_cuda_attention_project_batch8(ColiCudaTensor *kv_b,ColiCudaTensor *o_proj, + float *out,const float *q, + const uint8_t *latent,const float *lsc, + const uint8_t *rope,const float *rsc, + int S,int H,int Q,int R, + int V,int K,int T,float attention_scale); + COLI_CUDA_DLLEXPORT void coli_cuda_tensor_free(ColiCudaTensor *tensor); COLI_CUDA_DLLEXPORT size_t coli_cuda_tensor_bytes(const ColiCudaTensor *tensor); COLI_CUDA_DLLEXPORT int coli_cuda_tensor_device(const ColiCudaTensor *tensor); diff --git a/c/backend_loader.c b/c/backend_loader.c index b743d1b90..9ed91b520 100644 --- a/c/backend_loader.c +++ b/c/backend_loader.c @@ -63,6 +63,9 @@ typedef int (*fn_attention_absorb_kvdev)(ColiCudaTensor *kv_b,float *ctx,const f typedef int (*fn_attention_project_batch)(ColiCudaTensor *kv_b,ColiCudaTensor *o_proj, float *out,const float *q,const float *latent, const float *rope,int S,int H,int Q,int R, int V,int K,int T,float attention_scale); typedef int (*fn_attention_project_batch_dev)(ColiCudaTensor *kv_b,ColiCudaTensor *o_proj, float *out,const float *q_dev,const float *latent_dev,const float *rope_dev, int S,int H,int Q,int R,int V,int K,int T,float scale); typedef int (*fn_attention_project_batch_dev_out)(ColiCudaTensor *kv_b,ColiCudaTensor *o_proj, float *out_dev,const float *q_dev,const float *latent_dev,const float *rope_dev, int S,int H,int Q,int R,int V,int K,int T,float scale); +typedef int (*fn_attention_absorb8)(ColiCudaTensor *kv_b,float *ctx,const float *q, const uint8_t *latent,const float *lsc,const uint8_t *rope,const float *rsc, int H,int Q,int R,int V,int K,int T,float attention_scale); +typedef int (*fn_attention_absorb_batch8)(ColiCudaTensor *kv_b,float *ctx,const float *q, const uint8_t *latent,const float *lsc,const uint8_t *rope,const float *rsc, int S,int H,int Q,int R,int V,int K,int T,float attention_scale); +typedef int (*fn_attention_project_batch8)(ColiCudaTensor *kv_b,ColiCudaTensor *o_proj, float *out,const float *q,const uint8_t *latent,const float *lsc, const uint8_t *rope,const float *rsc,int S,int H,int Q,int R,int V,int K,int T,float attention_scale); typedef int (*fn_pipe_add)(int device,float *x_dev,const float *t_dev,size_t n); typedef void * (*fn_pipe_alloc)(int device,size_t bytes); typedef int (*fn_pipe_copy2d)(int device,float *dst,int dpitch,const float *src, int spitch,int width,int height); @@ -109,6 +112,9 @@ static struct { fn_attention_project_batch attention_project_batch; fn_attention_project_batch_dev attention_project_batch_dev; fn_attention_project_batch_dev_out attention_project_batch_dev_out; + fn_attention_absorb8 attention_absorb8; + fn_attention_absorb_batch8 attention_absorb_batch8; + fn_attention_project_batch8 attention_project_batch8; fn_pipe_add pipe_add; fn_pipe_alloc pipe_alloc; fn_pipe_copy2d pipe_copy2d; @@ -202,6 +208,9 @@ static int coli_cuda_load(void){ RESOLVE(attention_project_batch, fn_attention_project_batch) RESOLVE(attention_project_batch_dev, fn_attention_project_batch_dev) RESOLVE(attention_project_batch_dev_out, fn_attention_project_batch_dev_out) + RESOLVE(attention_absorb8, fn_attention_absorb8) + RESOLVE(attention_absorb_batch8, fn_attention_absorb_batch8) + RESOLVE(attention_project_batch8, fn_attention_project_batch8) RESOLVE(pipe_add, fn_pipe_add) RESOLVE(pipe_alloc, fn_pipe_alloc) RESOLVE(pipe_copy2d, fn_pipe_copy2d) @@ -327,6 +336,21 @@ int coli_cuda_attention_absorb_batch(ColiCudaTensor *kv_b,float *ctx,const float return g_cuda.attention_absorb_batch(kv_b, ctx, q, latent, rope, S, H, Q, R, V, K, T, attention_scale); } +int coli_cuda_attention_absorb8(ColiCudaTensor *kv_b,float *ctx,const float *q, const uint8_t *latent,const float *lsc,const uint8_t *rope,const float *rsc, int H,int Q,int R,int V,int K,int T,float attention_scale){ + if(!g_cuda.available){ return 0; } + return g_cuda.attention_absorb8(kv_b, ctx, q, latent, lsc, rope, rsc, H, Q, R, V, K, T, attention_scale); +} + +int coli_cuda_attention_absorb_batch8(ColiCudaTensor *kv_b,float *ctx,const float *q, const uint8_t *latent,const float *lsc,const uint8_t *rope,const float *rsc, int S,int H,int Q,int R,int V,int K,int T,float attention_scale){ + if(!g_cuda.available){ return 0; } + return g_cuda.attention_absorb_batch8(kv_b, ctx, q, latent, lsc, rope, rsc, S, H, Q, R, V, K, T, attention_scale); +} + +int coli_cuda_attention_project_batch8(ColiCudaTensor *kv_b,ColiCudaTensor *o_proj, float *out,const float *q,const uint8_t *latent,const float *lsc, const uint8_t *rope,const float *rsc,int S,int H,int Q,int R,int V,int K,int T,float attention_scale){ + if(!g_cuda.available){ return 0; } + return g_cuda.attention_project_batch8(kv_b, o_proj, out, q, latent, lsc, rope, rsc, S, H, Q, R, V, K, T, attention_scale); +} + int coli_cuda_attention_absorb_batch_dev(ColiCudaTensor *kv_b_shard,float *ctx_dev, const float *q_dev,const float *latent_dev,const float *rope_dev, int S,int H,int Q,int R,int V,int K,int T,float scale){ if(!g_cuda.available){ return 0; } return g_cuda.attention_absorb_batch_dev(kv_b_shard, ctx_dev, q_dev, latent_dev, rope_dev, S, H, Q, R, V, K, T, scale); diff --git a/c/decode_batch.h b/c/decode_batch.h index ad3d2edaa..8437ed955 100644 --- a/c/decode_batch.h +++ b/c/decode_batch.h @@ -2,6 +2,7 @@ #define COLIBRI_DECODE_BATCH_H #include +#include #include #include @@ -12,6 +13,12 @@ static inline float *coli_kv_row(float *base, int position, int width) return base + (size_t)position * (size_t)width; } +/* KV8 twin: same row arithmetic on the fp8 (e4m3) byte cache. */ +static inline uint8_t *coli_kv_row8(uint8_t *base, int position, int width) +{ + return base + (size_t)position * (size_t)width; +} + typedef struct { unsigned long long id, bytes; int slot, max_tokens; diff --git a/c/glm.c b/c/glm.c index 7e05ad944..3cbf2b866 100644 --- a/c/glm.c +++ b/c/glm.c @@ -49,6 +49,7 @@ #include "grammar.h" /* metodo F: draft grammaticali (#48) */ #include "schema_gbnf.h" /* SCHEMA=: JSON-Schema -> GBNF for method F */ #include "decode_batch.h" +#include "kv_fp8.h" /* KV8=1: cache latente in fp8 e4m3 + scala per-riga */ #ifdef _OPENMP #include /* scratch per-thread nell'attention */ #else @@ -146,6 +147,8 @@ typedef struct { int eid; QT g,u,d; uint8_t *slab; float *fslab; typedef struct { float **Lc, **Rc, **Ic; + uint8_t **Lc8, **Rc8; /* KV8: righe latenti fp8 e4m3 (Lc/Rc restano NULL) */ + float **Lsc, **Rsc; /* KV8: scala amax/448 per riga (per token, per layer) */ int *kv_start, max_t; int disk_nrec; char disk_path[2048]; @@ -168,6 +171,7 @@ typedef struct { * k_rot [qk_rope] (576 vs 32768 valori/token). k_nope e value si ricostruiscono al * volo con kv_b. E' cio' che rende gestibile il contesto su 15 GB (64 teste, no GQA). */ float **Lc, **Rc; int max_t; /* alias della KVState attiva */ + uint8_t **Lc8, **Rc8; float **Lsc, **Rsc; /* alias KV8 (fp8 + scale) della KVState attiva */ int *kv_start; /* prima pos valida nella KV del layer (MTP: parziale) */ KVState *kv; ESlot **ecache; int *ecn; int ecap; /* LRU expert per-layer */ @@ -2181,6 +2185,10 @@ static void qt_matvec_rows(const QT *t, int r0, int n, const float *x, float *y) } } static int g_absorb=-1; +/* KV8=1: cache latente Lc/Rc in fp8 e4m3 + scala f32 per riga (~4x meno RAM del f32). + * Coperti CPU e CUDA (kernel absorb8); auto-off sotto COLI_METAL, e forza + * COLI_CUDA_PIPE=0 (l'ombra KV su device e il pipe-prefill leggono righe f32). */ +static int g_kv8=0; #ifdef COLI_CUDA static int g_cuda_pipe=0; /* COLI_CUDA_PIPE=1: prefill attention chain resident on the layer home device */ #endif /* ABSORB: -1 auto (decode S<=4), 0 mai, 1 sempre (test) */ @@ -2405,16 +2413,26 @@ static void attention_rows(Model *m, Layer *l, int layer, float *x, int S, int p float *qfull=Q+(int64_t)s*H*qh; for(int h=0;hqk_nope, pos, c); const float *cs=comp+(int64_t)s*cw; - float *Ldst=coli_kv_row(ks->Lc[layer],pos,c->kv_lora); - float *Rdst=coli_kv_row(ks->Rc[layer],pos,c->qk_rope); #ifdef COLI_CUDA if(ks==m->kv&&m->kv_dev_valid&&layer<=c->n_layers&&m->kv_dev_valid[layer]>pos) m->kv_dev_valid[layer]=pos; /* riga riscritta: l'ombra si accorcia */ #endif - memcpy(Ldst, cs, c->kv_lora*sizeof(float)); - rmsnorm(Ldst, Ldst, l->kv_a_ln, c->kv_lora, c->eps); /* latente normato */ - memcpy(Rdst, cs+c->kv_lora, c->qk_rope*sizeof(float)); - rope_interleave(Rdst, pos, c); /* k_rot roped, condiviso fra teste */ + if(g_kv8){ + /* KV8: norma+rope sul residuo di comp (scratch, mai riletto), poi + * quantizza riga+scala. E' IL produttore caldo: ogni token, ogni layer. */ + float *Ls=comp+(int64_t)s*cw, *Rs=Ls+c->kv_lora; + rmsnorm(Ls, Ls, l->kv_a_ln, c->kv_lora, c->eps); /* latente normato */ + rope_interleave(Rs, pos, c); /* k_rot roped */ + ks->Lsc[layer][pos]=coli_kv8_quant_row(Ls, coli_kv_row8(ks->Lc8[layer],pos,c->kv_lora), c->kv_lora); + ks->Rsc[layer][pos]=coli_kv8_quant_row(Rs, coli_kv_row8(ks->Rc8[layer],pos,c->qk_rope), c->qk_rope); + } else { + float *Ldst=coli_kv_row(ks->Lc[layer],pos,c->kv_lora); + float *Rdst=coli_kv_row(ks->Rc[layer],pos,c->qk_rope); + memcpy(Ldst, cs, c->kv_lora*sizeof(float)); + rmsnorm(Ldst, Ldst, l->kv_a_ln, c->kv_lora, c->eps); /* latente normato */ + memcpy(Rdst, cs+c->kv_lora, c->qk_rope*sizeof(float)); + rope_interleave(Rdst, pos, c); /* k_rot roped, condiviso fra teste */ + } } /* ---- DSA lightning indexer ---- * Layer FULL: k_idx dei token nuovi in cache + selezione top-k per query (riusata @@ -2525,10 +2543,16 @@ static void attention_rows(Model *m, Layer *l, int layer, float *x, int S, int p qs+(int64_t)l->shard_h0[d]*S*qh+(int64_t)s*l->shard_hn[d]*qh, Q+((int64_t)s*H+l->shard_h0[d])*qh,(size_t)l->shard_hn[d]*qh*sizeof(float)); #pragma omp parallel for schedule(static) reduction(&:ok) - for(int d=0;dkv_b_shard[d], - cs+(int64_t)l->shard_h0[d]*S*vh,qs+(int64_t)l->shard_h0[d]*S*qh, - coli_kv_row(m->Lc[layer],st0,kvl),coli_kv_row(m->Rc[layer],st0,c->qk_rope), - S,l->shard_hn[d],c->qk_nope,c->qk_rope,vh,kvl,nt,c->attn_scale); + for(int d=0;dkv_b_shard[d], + cs+(int64_t)l->shard_h0[d]*S*vh,qs+(int64_t)l->shard_h0[d]*S*qh, + coli_kv_row8(m->Lc8[layer],st0,kvl),m->Lsc[layer]+st0, + coli_kv_row8(m->Rc8[layer],st0,c->qk_rope),m->Rsc[layer]+st0, + S,l->shard_hn[d],c->qk_nope,c->qk_rope,vh,kvl,nt,c->attn_scale) + :coli_cuda_attention_absorb_batch(l->kv_b_shard[d], + cs+(int64_t)l->shard_h0[d]*S*vh,qs+(int64_t)l->shard_h0[d]*S*qh, + coli_kv_row(m->Lc[layer],st0,kvl),coli_kv_row(m->Rc[layer],st0,c->qk_rope), + S,l->shard_hn[d],c->qk_nope,c->qk_rope,vh,kvl,nt,c->attn_scale); if(ok)for(int d=0;dshard_h0[d])*vh, cs+(int64_t)l->shard_h0[d]*S*vh+(int64_t)s*l->shard_hn[d]*vh, @@ -2537,9 +2561,14 @@ static void attention_rows(Model *m, Layer *l, int layer, float *x, int S, int p } else if(cuda_absorb&&l->kv_b.cuda_eligible&&l->o.cuda_eligible&& qt_cuda_upload(&l->kv_b)&&qt_cuda_upload(&l->o)){ int st0=m->kv_start[layer],nt=pos_base+S-st0; - cuda_core=cuda_projected=coli_cuda_attention_project_batch(l->kv_b.cuda,l->o.cuda,out,Q, - coli_kv_row(m->Lc[layer],st0,kvl),coli_kv_row(m->Rc[layer],st0,c->qk_rope), - S,H,c->qk_nope,c->qk_rope,vh,kvl,nt,c->attn_scale); + cuda_core=cuda_projected=g_kv8 + ?coli_cuda_attention_project_batch8(l->kv_b.cuda,l->o.cuda,out,Q, + coli_kv_row8(m->Lc8[layer],st0,kvl),m->Lsc[layer]+st0, + coli_kv_row8(m->Rc8[layer],st0,c->qk_rope),m->Rsc[layer]+st0, + S,H,c->qk_nope,c->qk_rope,vh,kvl,nt,c->attn_scale) + :coli_cuda_attention_project_batch(l->kv_b.cuda,l->o.cuda,out,Q, + coli_kv_row(m->Lc[layer],st0,kvl),coli_kv_row(m->Rc[layer],st0,c->qk_rope), + S,H,c->qk_nope,c->qk_rope,vh,kvl,nt,c->attn_scale); } else if(S<=4&&g_cuda_enabled&&getenv("COLI_CUDA_ATTN")&&atoi(getenv("COLI_CUDA_ATTN"))&& l->kv_b.cuda_eligible&&qt_cuda_upload(&l->kv_b)){ cuda_core=1; @@ -2554,10 +2583,15 @@ static void attention_rows(Model *m, Layer *l, int layer, float *x, int S, int p m->kv_dev_R[layer]+(size_t)st0*c->qk_rope,H,c->qk_nope,c->qk_rope, vh,kvl,nt,c->attn_scale); if(!cuda_core) - cuda_core=coli_cuda_attention_absorb(l->kv_b.cuda,ctx+(int64_t)s*H*vh, - Q+(int64_t)s*H*qh,coli_kv_row(ks->Lc[layer],st0,kvl), - coli_kv_row(ks->Rc[layer],st0,c->qk_rope),H,c->qk_nope,c->qk_rope, - vh,kvl,nt,c->attn_scale); + cuda_core=g_kv8 + ?coli_cuda_attention_absorb8(l->kv_b.cuda,ctx+(int64_t)s*H*vh, + Q+(int64_t)s*H*qh,coli_kv_row8(ks->Lc8[layer],st0,kvl), + ks->Lsc[layer]+st0,coli_kv_row8(ks->Rc8[layer],st0,c->qk_rope), + ks->Rsc[layer]+st0,H,c->qk_nope,c->qk_rope,vh,kvl,nt,c->attn_scale) + :coli_cuda_attention_absorb(l->kv_b.cuda,ctx+(int64_t)s*H*vh, + Q+(int64_t)s*H*qh,coli_kv_row(ks->Lc[layer],st0,kvl), + coli_kv_row(ks->Rc[layer],st0,c->qk_rope),H,c->qk_nope,c->qk_rope, + vh,kvl,nt,c->attn_scale); } } #endif @@ -2577,17 +2611,35 @@ static void attention_rows(Model *m, Layer *l, int layer, float *x, int S, int p const int *tlist = ns ? dsel+(int64_t)s*dtopk : NULL; int nt = ns ? ns : pos+1-st0; for(int jj=0;jjLc[layer],t,kvl); - const float *kr=coli_kv_row(ks->Rc[layer],t,c->qk_rope); - float a=0; for(int i=0;iqk_rope;d++) a+=qr[d]*kr[d]; + float a=0; + if(g_kv8){ + /* LUT-dequant inline nel dot; la scala per-riga esce dalla + * somma: score = Lsc·Σ q·lut[b] + Rsc·Σ qr·lut[b] */ + const uint8_t *Lt=coli_kv_row8(ks->Lc8[layer],t,kvl); + const uint8_t *kr=coli_kv_row8(ks->Rc8[layer],t,c->qk_rope); + float al=0, ar=0; + for(int i=0;iqk_rope;d++) ar+=qr[d]*coli_fp8_lut[kr[d]]; + a=al*ks->Lsc[layer][t]+ar*ks->Rsc[layer][t]; + } else { + const float *Lt=coli_kv_row(ks->Lc[layer],t,kvl); + const float *kr=coli_kv_row(ks->Rc[layer],t,c->qk_rope); + for(int i=0;iqk_rope;d++) a+=qr[d]*kr[d]; + } sc[jj]=a*c->attn_scale; } softmax(sc,nt); float clat[512]; memset(clat,0,kvl*sizeof(float)); for(int jj=0;jjLc[layer],t,kvl); - float a=sc[jj]; for(int i=0;iLc8[layer],t,kvl); + float a=sc[jj]*ks->Lsc[layer][t]; /* la scala si fonde nel peso */ + for(int i=0;iLc[layer],t,kvl); + float a=sc[jj]; for(int i=0;ikv_b, rbase+r0v, vh, clat, ctx+((int64_t)s*H+h)*vh); } } @@ -2601,7 +2653,17 @@ static void attention_rows(Model *m, Layer *l, int layer, float *x, int S, int p m->t_aproj+=now_s()-ta0; double tk0=now_s(); int stL=m->kv_start[layer]; float *kvb_all=falloc((int64_t)Tk*kvb_dim); - matmul_qt(kvb_all+(int64_t)stL*kvb_dim, m->Lc[layer]+(int64_t)stL*c->kv_lora, &l->kv_b, Tk-stL); + if(g_kv8){ + /* staging f32 del latente dequantizzato: kv_b vuole righe float. Il buffer + * [Tk-stL,kvl] e' rumore rispetto a kvb_all [Tk,H*(nope+vh)] gia' allocato. */ + float *Lf=falloc((int64_t)(Tk-stL)*c->kv_lora); + for(int t=stL;tLc8[layer],t,c->kv_lora), m->Lsc[layer][t], + Lf+(int64_t)(t-stL)*c->kv_lora, c->kv_lora); + matmul_qt(kvb_all+(int64_t)stL*kvb_dim, Lf, &l->kv_b, Tk-stL); + free(Lf); + } else + matmul_qt(kvb_all+(int64_t)stL*kvb_dim, m->Lc[layer]+(int64_t)stL*c->kv_lora, &l->kv_b, Tk-stL); m->t_kvb += now_s()-tk0; /* 3) attenzione causale: score = q_pass·k_nope + q_rot·k_rot * (punteggi sul heap, per-thread: vedi il commento nel ramo absorb) */ @@ -2620,9 +2682,15 @@ static void attention_rows(Model *m, Layer *l, int layer, float *x, int S, int p int nt = ns ? ns : pos+1-st0; for(int jj=0;jjqk_nope+vh); - const float *kr=m->Rc[layer]+(int64_t)t*c->qk_rope; float a=0; for(int d=0;dqk_nope;d++) a+=qp[d]*kn[d]; - for(int d=0;dqk_rope;d++) a+=qr[d]*kr[d]; + if(g_kv8){ + const uint8_t *kr=coli_kv_row8(m->Rc8[layer],t,c->qk_rope); + float ar=0; for(int d=0;dqk_rope;d++) ar+=qr[d]*coli_fp8_lut[kr[d]]; + a+=ar*m->Rsc[layer][t]; + } else { + const float *kr=m->Rc[layer]+(int64_t)t*c->qk_rope; + for(int d=0;dqk_rope;d++) a+=qr[d]*kr[d]; + } sc[jj]=a*c->attn_scale; } softmax(sc,nt); @@ -3793,6 +3861,10 @@ static void kv_alloc(Model *m, int max_t){ if(g_metal_enabled){ coli_metal_unregister(k->Lc[i]); coli_metal_unregister(k->Rc[i]); } #endif free(k->Lc[i]); free(k->Rc[i]); } free(k->Lc); free(k->Rc); } + if(k->Lc8){ for(int i=0;in_layers+1;i++){ free(k->Lc8[i]); free(k->Rc8[i]); + free(k->Lsc[i]); free(k->Rsc[i]); } + free(k->Lc8); free(k->Rc8); free(k->Lsc); free(k->Rsc); + k->Lc8=k->Rc8=NULL; k->Lsc=k->Rsc=NULL; } if(k->Ic){ for(int i=0;in_layers;i++) free(k->Ic[i]); free(k->Ic); k->Ic=NULL; } if(m->has_dsa){ k->Ic=calloc(c->n_layers,sizeof(float*)); @@ -3801,6 +3873,19 @@ static void kv_alloc(Model *m, int max_t){ k->max_t=max_t; int NR=c->n_layers+1; /* riga extra: KV del layer MTP */ k->Lc=calloc(NR,sizeof(float*)); k->Rc=calloc(NR,sizeof(float*)); + if(g_kv8){ + /* KV8: byte fp8 (non float) + una scala f32 per riga; Lc/Rc restano NULL. + * (576+8)/2304 B/token/layer contro i 2304 del f32: ~3.9x di RAM in meno. */ + coli_fp8_lut_init(); + k->Lc8=calloc(NR,sizeof(uint8_t*)); k->Rc8=calloc(NR,sizeof(uint8_t*)); + k->Lsc=calloc(NR,sizeof(float*)); k->Rsc=calloc(NR,sizeof(float*)); + for(int i=0;iLc8[i]=malloc((size_t)max_t*c->kv_lora); + k->Rc8[i]=malloc((size_t)max_t*c->qk_rope); + k->Lsc[i]=falloc(max_t); k->Rsc[i]=falloc(max_t); + if(!k->Lc8[i]||!k->Rc8[i]){fprintf(stderr,"OOM kv8\n");exit(1);} + } + } else for(int i=0;iLc[i]=falloc((int64_t)max_t*c->kv_lora); k->Rc[i]=falloc((int64_t)max_t*c->qk_rope); #ifdef COLI_METAL @@ -3817,12 +3902,14 @@ static void kv_alloc(Model *m, int max_t){ #endif } m->Lc=k->Lc; m->Rc=k->Rc; m->Ic=k->Ic; m->max_t=k->max_t; m->kv_start=k->kv_start; + m->Lc8=k->Lc8; m->Rc8=k->Rc8; m->Lsc=k->Lsc; m->Rsc=k->Rsc; } static void kv_bind(Model *m, KVState *k){ if(m->kv!=k && m->kv_dev_valid) /* ombra legata al KVState corrente */ for(int i=0;ic.n_layers+1;i++) m->kv_dev_valid[i]=0; m->kv=k; m->Lc=k->Lc; m->Rc=k->Rc; m->Ic=k->Ic; + m->Lc8=k->Lc8; m->Rc8=k->Rc8; m->Lsc=k->Lsc; m->Rsc=k->Rsc; m->max_t=k->max_t; m->kv_start=k->kv_start; } @@ -3870,7 +3957,10 @@ static float *step_decode_batch(Model *m, const DecodeRow *rows, int S){ free(x); return NULL; } for(int l=0;ln_layers;l++){ - if(!rows[s].kv->Lc[l] || !rows[s].kv->Rc[l] || + if((g_kv8 ? (!rows[s].kv->Lc8 || !rows[s].kv->Rc8 || + !rows[s].kv->Lc8[l] || !rows[s].kv->Rc8[l] || + !rows[s].kv->Lsc[l] || !rows[s].kv->Rsc[l]) + : (!rows[s].kv->Lc[l] || !rows[s].kv->Rc[l])) || rows[s].kv->kv_start[l]<0 || rows[s].kv->kv_start[l]>rows[s].pos || (m->has_dsa && c->idx_type[l] && (!rows[s].kv->Ic || !rows[s].kv->Ic[l]))){ free(x); return NULL; } @@ -4629,33 +4719,43 @@ static void repin_pass_limit(Model *m,int limit){ * append lascia nrec vecchio = file coerente. La riga KV del layer MTP non si salva: * al resume kv_start=-1 e la finestra di draft riparte da sola. */ static int g_kvsave=1; -#define KV_MAGIC "COLIKV1\0" +#define KV_MAGIC "COLIKV1\0" /* v1: righe Lc/Rc f32 */ +#define KV_MAGIC2 "COLIKV2\0" /* v2 (KV8): righe fp8 e4m3 + scala f32 per riga */ static void kv_hdr(Model *m, int32_t *h, int nrec){ Cfg *c=&m->c; int nic=0; for(int i=0;in_layers;i++) if(m->Ic && m->Ic[i]) nic++; h[0]=c->n_layers; h[1]=c->kv_lora; h[2]=c->qk_rope; - h[3]=m->has_dsa?c->index_hd:0; h[4]=nic; h[5]=c->vocab; h[6]=nrec; h[7]=0; + h[3]=m->has_dsa?c->index_hd:0; h[4]=nic; h[5]=c->vocab; h[6]=nrec; h[7]=g_kv8?1:0; } /* Bytes of one on-disk record: [tok i32][Lc+Rc per layer][Ic per DSA layer]. - * Layout matches what kv_disk_append writes and kv_disk_load reads. */ + * Layout matches what kv_disk_append writes and kv_disk_load reads. Under KV8 + * (v2 format) Lc/Rc are fp8 bytes + one f32 scale each: ~46 KB/token instead + * of ~182 (4x less write-amplification on ZFS). */ static int64_t kv_rec_bytes(Model *m){ Cfg *c=&m->c; - int64_t rec = 4 + (int64_t)c->n_layers*(c->kv_lora+c->qk_rope)*4; + int64_t rec = 4 + (g_kv8 ? (int64_t)c->n_layers*(c->kv_lora+c->qk_rope+8) + : (int64_t)c->n_layers*(c->kv_lora+c->qk_rope)*4); if(m->has_dsa) for(int i=0;in_layers;i++) if(m->Ic[i]) rec+=(int64_t)c->index_hd*4; return rec; } /* Open the persistent handle lazily; write the header if the file is new. After * this returns successfully, k->disk_fp is valid for the engine's lifetime and - * positioned at end-of-header (nrec==0 case) or wherever the caller seeks. */ + * positioned at end-of-header (nrec==0 case) or wherever the caller seeks. + * Formato del file != formato attivo (KV8 acceso o spento tra un run e l'altro): + * il file si riscrive da zero — e' il punto dove l'upgrade v1->v2 diventa reale. */ static int kv_disk_open(Model *m){ KVState *k=m->kv; if(k->disk_fp) return 1; k->disk_fp=fopen(k->disk_path,"r+b"); - if(!k->disk_fp){ /* not there yet -> create + header */ + if(k->disk_fp){ char mg[8]; + if(fread(mg,1,8,k->disk_fp)!=8 || memcmp(mg,g_kv8?KV_MAGIC2:KV_MAGIC,8)){ + fclose(k->disk_fp); k->disk_fp=NULL; k->disk_nrec=0; } + } + if(!k->disk_fp){ /* not there yet (or wrong format) -> create + header */ k->disk_fp=fopen(k->disk_path,"wb"); if(!k->disk_fp) return 0; int32_t h[8]; kv_hdr(m,h,0); - fwrite(KV_MAGIC,1,8,k->disk_fp); fwrite(h,4,8,k->disk_fp); + fwrite(g_kv8?KV_MAGIC2:KV_MAGIC,1,8,k->disk_fp); fwrite(h,4,8,k->disk_fp); fflush(k->disk_fp); fclose(k->disk_fp); k->disk_fp=fopen(k->disk_path,"r+b"); /* reopen r+b for append */ @@ -4692,8 +4792,15 @@ static void kv_disk_append(Model *m, const int *hist, int len){ uint8_t *b=k->disk_buf; /* pack token + every layer into one record */ *(int32_t*)b = hist[p]; b+=4; for(int i=0;in_layers;i++){ - memcpy(b, m->Lc[i]+(int64_t)p*c->kv_lora, (size_t)c->kv_lora*4); b+=c->kv_lora*4; - memcpy(b, m->Rc[i]+(int64_t)p*c->qk_rope,(size_t)c->qk_rope*4); b+=c->qk_rope*4; + if(g_kv8){ /* v2: fp8 + scala per riga, stesso staging */ + memcpy(b, coli_kv_row8(m->Lc8[i],p,c->kv_lora), (size_t)c->kv_lora); b+=c->kv_lora; + memcpy(b, &m->Lsc[i][p], 4); b+=4; + memcpy(b, coli_kv_row8(m->Rc8[i],p,c->qk_rope), (size_t)c->qk_rope); b+=c->qk_rope; + memcpy(b, &m->Rsc[i][p], 4); b+=4; + } else { + memcpy(b, m->Lc[i]+(int64_t)p*c->kv_lora, (size_t)c->kv_lora*4); b+=c->kv_lora*4; + memcpy(b, m->Rc[i]+(int64_t)p*c->qk_rope,(size_t)c->qk_rope*4); b+=c->qk_rope*4; + } } if(m->has_dsa) for(int i=0;in_layers;i++) if(m->Ic[i]){ memcpy(b, m->Ic[i]+(int64_t)p*c->index_hd, (size_t)c->index_hd*4); b+=c->index_hd*4; @@ -4711,30 +4818,58 @@ static int kv_disk_load(Model *m, int *hist, int maxctx){ Cfg *c=&m->c; FILE *f=fopen(k->disk_path,"rb"); if(!f) return 0; char mg[8]; int32_t h[8], w[8]; kv_hdr(m,w,0); - if(fread(mg,1,8,f)!=8 || memcmp(mg,KV_MAGIC,8) || fread(h,4,8,f)!=8 || + int dt=-1; /* dtype del FILE: 0=f32 (v1), 1=fp8 (v2) */ + if(fread(mg,1,8,f)==8){ + if(!memcmp(mg,KV_MAGIC,8)) dt=0; else if(!memcmp(mg,KV_MAGIC2,8)) dt=1; } + if(dt<0 || fread(h,4,8,f)!=8 || h[0]!=w[0]||h[1]!=w[1]||h[2]!=w[2]||h[3]!=w[3]||h[4]!=w[4]||h[5]!=w[5]){ fprintf(stderr,"[KV] ignoring .coli_kv from a different model or version\n"); fclose(f); return 0; } + if(dt==1 && !g_kv8){ + fprintf(stderr,"[KV] .coli_kv is fp8 (saved under KV8=1): starting over (set KV8=1 to resume it)\n"); + fclose(f); return 0; } int nrec=h[6]; if(nrec<1){ fclose(f); return 0; } if(nrec>=maxctx-8-g_draft){ fprintf(stderr,"[KV] saved conversation (%d tokens) exceeds the context: starting over\n",nrec); fclose(f); return 0; } double t0=now_s(); + /* v1 sotto KV8: righe f32 lette in staging e quantizzate al volo */ + float *stage = (g_kv8&&dt==0) ? falloc(c->kv_lora>c->qk_rope?c->kv_lora:c->qk_rope) : NULL; for(int p=0;pn_layers;i++){ - if(fread(m->Lc[i]+(int64_t)p*c->kv_lora, 4, c->kv_lora, f)!=(size_t)c->kv_lora || - fread(m->Rc[i]+(int64_t)p*c->qk_rope, 4, c->qk_rope, f)!=(size_t)c->qk_rope){ nrec=p; goto out; } + if(dt==1){ /* v2: fp8+scala, gia' nel formato in RAM */ + if(fread(coli_kv_row8(m->Lc8[i],p,c->kv_lora), 1, c->kv_lora, f)!=(size_t)c->kv_lora || + fread(&m->Lsc[i][p], 4, 1, f)!=1 || + fread(coli_kv_row8(m->Rc8[i],p,c->qk_rope), 1, c->qk_rope, f)!=(size_t)c->qk_rope || + fread(&m->Rsc[i][p], 4, 1, f)!=1){ nrec=p; goto out; } + } else if(g_kv8){ + if(fread(stage, 4, c->kv_lora, f)!=(size_t)c->kv_lora){ nrec=p; goto out; } + m->Lsc[i][p]=coli_kv8_quant_row(stage, coli_kv_row8(m->Lc8[i],p,c->kv_lora), c->kv_lora); + if(fread(stage, 4, c->qk_rope, f)!=(size_t)c->qk_rope){ nrec=p; goto out; } + m->Rsc[i][p]=coli_kv8_quant_row(stage, coli_kv_row8(m->Rc8[i],p,c->qk_rope), c->qk_rope); + } else { + if(fread(m->Lc[i]+(int64_t)p*c->kv_lora, 4, c->kv_lora, f)!=(size_t)c->kv_lora || + fread(m->Rc[i]+(int64_t)p*c->qk_rope, 4, c->qk_rope, f)!=(size_t)c->qk_rope){ nrec=p; goto out; } + } } if(m->has_dsa) for(int i=0;in_layers;i++) if(m->Ic[i]) if(fread(m->Ic[i]+(int64_t)p*c->index_hd, 4, c->index_hd, f)!=(size_t)c->index_hd){ nrec=p; goto out; } } out: - fclose(f); + fclose(f); free(stage); if(nrec>0){ if(m->has_mtp) m->kv_start[c->n_layers]=-1; /* la finestra MTP riparte da sola */ fprintf(stderr,"[KV] resumed conversation from disk: %d tokens in %.1fs (no re-prefill)\n", nrec, now_s()-t0); + if(g_kv8 && dt==0){ + /* upgrade v1->v2: il file f32 non puo' ricevere append fp8. disk_nrec=0 e + * il file resta INTATTO (un crash prima del primo save non perde nulla): + * al primo append il controllo magic lo riscrive da zero in v2. */ + k->disk_nrec=0; + fprintf(stderr,"[KV] f32 .coli_kv quantized in RAM; will be rewritten as fp8 (v2) at next save\n"); + return nrec; + } } k->disk_nrec=nrec; return nrec; @@ -4760,8 +4895,11 @@ static void serve_ctx_free(Model *m, ServeCtx *s){ if(k->disk_fp){ fclose(k->disk_fp); k->disk_fp=NULL; } free(k->disk_buf); k->disk_buf=NULL; if(k->Lc) for(int i=0;iLc[i]); free(k->Rc[i]); } + if(k->Lc8) for(int i=0;iLc8[i]); free(k->Rc8[i]); + free(k->Lsc[i]); free(k->Rsc[i]); } if(k->Ic) for(int i=0;ic.n_layers;i++) free(k->Ic[i]); - free(k->Lc); free(k->Rc); free(k->Ic); free(k->kv_start); free(s->hist); + free(k->Lc); free(k->Rc); free(k->Lc8); free(k->Rc8); free(k->Lsc); free(k->Rsc); + free(k->Ic); free(k->kv_start); free(s->hist); } typedef struct { @@ -4943,7 +5081,7 @@ static void run_serve_mux(Model *m, const char *snap){ } usage_save(m); for(int i=0;ikv=NULL; m->Lc=m->Rc=m->Ic=NULL; m->kv_start=NULL; m->max_t=0; + m->kv=NULL; m->Lc=m->Rc=m->Ic=NULL; m->Lc8=m->Rc8=NULL; m->Lsc=m->Rsc=NULL; m->kv_start=NULL; m->max_t=0; } static void run_serve(Model *m, const char *snap){ @@ -5113,7 +5251,7 @@ static void run_serve(Model *m, const char *snap){ #undef len #undef first for(int i=0;ikv=NULL; m->Lc=m->Rc=m->Ic=NULL; m->kv_start=NULL; m->max_t=0; + free(ctx); m->kv=NULL; m->Lc=m->Rc=m->Ic=NULL; m->Lc8=m->Rc8=NULL; m->Lsc=m->Rsc=NULL; m->kv_start=NULL; m->max_t=0; } static int *read_arr(jval*o,const char*k,int*n){ @@ -5578,7 +5716,12 @@ static int kv_slot_count(void){ } static double kv_pool_bytes(Model *m, int max_ctx){ - Cfg *c=&m->c; double one=(double)(c->n_layers+1)*max_ctx*(c->kv_lora+c->qk_rope)*4.0; + /* KV8: 1 byte/valore + 8 B di scale per token/layer, non 4 B/valore. E' questo + * conto che governa expert_avail e cap_for_ram: con KV8 il clamp PIN recupera + * ~35 GB/slot a 256k e KV_SLOTS=2 smette di demolire gli expert su disco. */ + Cfg *c=&m->c; + double one=(double)(c->n_layers+1)*max_ctx* + (g_kv8 ? (double)(c->kv_lora+c->qk_rope)+8.0 : (c->kv_lora+c->qk_rope)*4.0); if(m->has_dsa) for(int i=0;in_layers;i++) if(c->idx_type[i]) one+=(double)max_ctx*c->index_hd*4.0; int slots=kv_slot_count(); if(slots<1||slots>16) slots=1; @@ -5884,6 +6027,28 @@ int main(int argc, char **argv){ return 2; } #endif + /* KV8=1: KV-cache latente in fp8 e4m3. CPU e CUDA (kernel absorb8) sono coperti; + * lo shader Metal e il resident pipeline (COLI_CUDA_PIPE) leggono righe f32: + * finche' non esistono le varianti fp8, quei percorsi si spengono da soli. */ + g_kv8 = getenv("KV8")?atoi(getenv("KV8")):0; + if(g_kv8){ +#ifdef COLI_METAL + if(g_kv8 && g_metal_enabled){ + fprintf(stderr,"[KV8] COLI_METAL=1: fp8 KV unsupported on the Metal path; KV8 disabled\n"); + g_kv8=0; + } +#endif +#ifdef COLI_CUDA + if(g_kv8 && g_cuda_pipe){ + fprintf(stderr,"[KV8] COLI_CUDA_PIPE reads f32 KV rows; pipe disabled under KV8\n"); + g_cuda_pipe=0; + } +#endif + if(g_kv8){ + coli_fp8_lut_init(); + fprintf(stderr,"[KV8] latent KV cache in fp8 e4m3 + per-row scale (~3.9x less KV RAM)\n"); + } + } printf("== GLM C engine (glm_moe_dsa), cache=%d experts/layer | experts@%d-bit dense@%d-bit | idot: " IDOT_KERNEL " ==\n", cap, ebits, dbits); g_mem_avail_boot = mem_available_gb(); Model m; double t0=now_s(); model_init(&m,snap,cap,ebits,dbits); diff --git a/c/kv_fp8.h b/c/kv_fp8.h new file mode 100644 index 000000000..fc9aeeca6 --- /dev/null +++ b/c/kv_fp8.h @@ -0,0 +1,65 @@ +#ifndef COLIBRI_KV_FP8_H +#define COLIBRI_KV_FP8_H + +#include +#include +#include + +/* FP8 e4m3 (OCP "fn" variant: no infinities, max ±448, NaN = S.1111.111) per la + * KV-cache latente MLA. Il latente compresso e' sia key sia value nell'attention + * assorbita, quindi l'errore di quantizzazione entra negli score E nel context — + * lo stesso regime della pratica FP8-KV di DeepSeek-V3. Scala per-riga (amax/448, + * f32, Lc e Rc separate): dentro una riga la dinamica e' piccola, tra token no. + * EN: e4m3 storage for the MLA latent KV rows, per-row f32 amax scale. Decode is + * lut[byte]*scale; encode is bit-math + round-to-nearest-even, no libraries. */ + +static float coli_fp8_lut[256]; + +static void coli_fp8_lut_init(void){ + if(coli_fp8_lut[1]!=0.f) return; /* idempotente */ + for(int b=0;b<256;b++){ + int E=(b>>3)&0xF, M=b&7; + float v = E ? ldexpf(1.f+(float)M/8.f, E-7) /* normale: (1+M/8)*2^(E-7) */ + : ldexpf((float)M, -9); /* denormale: M/8 * 2^-6 */ + if(E==15 && M==7) v=0.f; /* codice NaN: mai scritto; inerte in lettura */ + coli_fp8_lut[b] = (b&0x80) ? -v : v; + } +} + +/* float -> e4m3, round-to-nearest-even, saturazione a ±448 (come __nv_cvt SATFINITE). + * NaN -> 0: una riga f32 con NaN avvelenerebbe comunque tutto; 0 e' inerte nei dot. */ +static inline uint8_t coli_fp8_enc(float f){ + union { float f; uint32_t u; } v; v.f=f; + uint8_t s=(uint8_t)((v.u>>24)&0x80u); + float a=fabsf(f); + if(!(a<=448.f)) return (a!=a) ? 0 : (uint8_t)(s|0x7e); /* NaN -> 0, inf/overflow -> ±448 */ + if(a<0x1p-6f){ /* griglia denormale: passo 2^-9 */ + int k=(int)rintf(a*0x1p9f); /* RNE su [0..8] */ + return (uint8_t)(s|(uint8_t)k); /* k==8 -> 0x08 = primo normale 2^-6 */ + } + int e; frexpf(a,&e); /* a = m*2^e, m in [0.5,1) */ + int E=e+6; /* esponente biased e4m3 (1..15) */ + int k=(int)rintf(ldexpf(a,3-(e-1))); /* mantissa RNE in [8,16] */ + if(k==16){ k=8; E++; } /* overflow di mantissa -> esponente su */ + if(E>15||(E==15&&k>14)) return (uint8_t)(s|0x7e); + return (uint8_t)(s|(uint8_t)(E<<3)|(uint8_t)(k-8)); +} + +/* quantizza una riga latente: scala amax/448 per-riga, ritorna la scala. + * La riga si decodifica come coli_fp8_lut[b]*scale. Riga tutta zero, subnormale + * (448/amax andrebbe a +inf) o non finita: byte 0 e scala 1, niente 1/0. */ +static inline float coli_kv8_quant_row(const float *src, uint8_t *dst, int n){ + float amax=0; + for(int i=0;iamax) amax=a; } + if(!(amax>1e-35f) || amax>3.4e38f){ memset(dst,0,(size_t)n); return 1.f; } + float inv=448.f/amax; + for(int i=0;i #include @@ -107,6 +108,45 @@ int main(int argc, char **argv) { !close_enough(actx,aref,2))return 1; coli_cuda_tensor_free(at); + /* KV8: same absorb case with e4m3-quantized latent/rope + per-row scales. + The reference is computed on the host from the DEQUANTIZED rows (exactly + what the kernel sees), so only accumulation order separates the two. */ + { + coli_fp8_lut_init(); + const float aw8[16]={1,0,0,0, 0,1,0,0, 0,0,1,0, 0,0,0,1}; + const float aq8[4]={1,2,.5f,-.5f}; + float al8f[12],ar8f[6]; + for(int i=0;i<12;i++)al8f[i]=std::sin((float)(i+1)*0.83f)*3.f; + for(int i=0;i<6;i++)ar8f[i]=std::cos((float)(i+1)*0.51f)*2.f; + uint8_t alq[12],arq[6];float alsc[3],arsc[3],ald[12],ard[6]; + for(int t=0;t<3;t++){ + alsc[t]=coli_kv8_quant_row(al8f+t*4,alq+t*4,4); + arsc[t]=coli_kv8_quant_row(ar8f+t*2,arq+t*2,2); + coli_kv8_dequant_row(alq+t*4,alsc[t],ald+t*4,4); + coli_kv8_dequant_row(arq+t*2,arsc[t],ard+t*2,2); + } + float sc8[3],ref8[2],got8[2]; + for(int t=0;t<3;t++)sc8[t]=aq8[0]*ald[t*4]+aq8[1]*ald[t*4+1]+aq8[2]*ard[t*2]+aq8[3]*ard[t*2+1]; + float m8=sc8[0],z8=0;for(int t=1;t<3;t++)m8=sc8[t]>m8?sc8[t]:m8; + for(int t=0;t<3;t++){sc8[t]=std::exp(sc8[t]-m8);z8+=sc8[t];}for(int t=0;t<3;t++)sc8[t]/=z8; + for(int v=0;v<2;v++){ref8[v]=0;for(int t=0;t<3;t++)ref8[v]+=sc8[t]*ald[t*4+2+v];} + ColiCudaTensor *at8=nullptr;if(!coli_cuda_tensor_upload(&at8,aw8,nullptr,0,4,4,d0))return 1; + if(!coli_cuda_attention_absorb8(at8,got8,aq8,alq,alsc,arq,arsc,1,2,2,2,4,3,1.f)|| + !close_enough(got8,ref8,2)){std::fprintf(stderr,"attention_absorb8 mismatch\n");return 1;} + /* batch twin, S=2 (query s attends T-S+s+1 rows): reference per query */ + float bq8[2*4]={1,2,.5f,-.5f, -1,.5f,1,2},bref[4],bgot[4]; + for(int s=0;s<2;s++){ + int ntk=3-2+s+1;float bs[3]; + for(int t=0;tbm?bs[t]:bm; + for(int t=0;t fp8 transition frees the f32 rows and allocates the byte + * cache + per-row scales; a second KV8 kv_alloc re-runs the fp8 free path; + * switching back must survive too (the KV8 arrays are freed and nulled). */ + g_kv8=1; + kv_alloc(&m,16); + for(int i=0;i f32 transition */ + if(m.kv->Lc8||m.kv->Rc8||m.kv->Lsc||m.kv->Rsc){ + fprintf(stderr,"KV8 arrays must be freed and nulled on the way back\n"); return 1; } + for(int i=0;iv2 + * upgrade path (quantize-on-load; the v1 file survives until the first save rewrites it + * as v2), the v2-under-f32 reject, and the format-mismatch self-heal in + * kv_disk_append (magic differs -> file rewritten from record 0). No model + * file needed: the disk paths only read c->n_layers/kv_lora/qk_rope/vocab. */ +#define main coli_glm_main_unused +#include "../glm.c" +#undef main + +#define PATH "tests/.coli_kv_test.tmp" +static int fails=0; +#define CHECK(cond, ...) do{ if(!(cond)){ fails++; \ + fprintf(stderr,"FAIL %s:%d: ",__FILE__,__LINE__); fprintf(stderr,__VA_ARGS__); fputc('\n',stderr);} }while(0) + +static float fill(int i,int p,int j){ return sinf(i*7.f+p*1.3f+j*0.37f)*(1.f+p); } + +int main(void){ + static Model m; + m.c.n_layers=2; m.c.kv_lora=8; m.c.qk_rope=4; m.c.vocab=256; + m.kv=calloc(1,sizeof(KVState)); + snprintf(m.kv->disk_path,sizeof(m.kv->disk_path),"%s",PATH); + remove(PATH); g_kvsave=1; g_draft=0; + int NP=5, hist[5]={11,22,33,44,55}, hist2[16]={0}; + + /* ---- v1: f32 append + load must round-trip exactly ---- */ + g_kv8=0; kv_alloc(&m,16); + for(int i=0;idisk_nrec==NP, "v1 append nrec=%d", m.kv->disk_nrec); + kv_alloc(&m,16); m.kv->disk_nrec=0; /* cache azzerata, si ricarica da disco */ + CHECK(kv_disk_load(&m,hist2,16)==NP, "v1 load"); + for(int p=0;pdisk_fp){ fclose(m.kv->disk_fp); m.kv->disk_fp=NULL; } /* "riavvio" del processo */ + g_kv8=1; kv_alloc(&m,16); m.kv->disk_nrec=0; + CHECK(kv_disk_load(&m,hist2,16)==NP, "v1->kv8 load"); + CHECK(m.kv->disk_nrec==0, "v1->kv8 leaves disk_nrec=0 for the full rewrite"); + { FILE *f=fopen(PATH,"rb"); char mg[8]={0}; + CHECK(f && fread(mg,1,8,f)==8 && !memcmp(mg,KV_MAGIC,8), + "v1 file must survive the upgrade load untouched"); + if(f) fclose(f); } + for(int i=0;ikv8 Lc[%d] p%d j%d: %g vs %g", i,p,j,got,want); + } + } + + /* ---- v2: fp8 append + load must round-trip BYTE-identical ---- */ + kv_disk_append(&m,hist,NP); + CHECK(m.kv->disk_nrec==NP, "v2 append nrec=%d", m.kv->disk_nrec); + { FILE *f=fopen(PATH,"rb"); char mg[8]={0}; + CHECK(f && fread(mg,1,8,f)==8 && !memcmp(mg,KV_MAGIC2,8), "v2 magic on disk"); + if(f) fclose(f); } + uint8_t keepL[2][5*8]; float keepS[2][5]; + for(int i=0;idisk_nrec=0; + CHECK(kv_disk_load(&m,hist2,16)==NP, "v2 load"); + for(int i=0;idisk_fp){ fclose(m.kv->disk_fp); m.kv->disk_fp=NULL; } /* "riavvio" del processo */ + g_kv8=0; kv_alloc(&m,16); m.kv->disk_nrec=0; + CHECK(kv_disk_load(&m,hist2,16)==0, "v2 under f32 must be rejected"); + + /* ---- append con formato diverso dal file: riscrittura dal record 0 ---- */ + for(int i=0;idisk_fp){ fclose(m.kv->disk_fp); m.kv->disk_fp=NULL; } /* "riavvio" del processo */ + m.kv->disk_nrec=3; /* stantio: la magic v2 lo invalida */ + kv_disk_append(&m,hist,NP); + { FILE *f=fopen(PATH,"rb"); char mg[8]={0}; + CHECK(f && fread(mg,1,8,f)==8 && !memcmp(mg,KV_MAGIC,8), "self-heal rewrote v1 magic"); + if(f) fclose(f); } + kv_alloc(&m,16); m.kv->disk_nrec=0; + CHECK(kv_disk_load(&m,hist2,16)==NP, "self-healed v1 load"); + for(int p=0;p +#include +#include +#include "../kv_fp8.h" +#include "../decode_batch.h" + +static int fails=0; +#define CHECK(cond, ...) do{ if(!(cond)){ fails++; \ + fprintf(stderr,"FAIL %s:%d: ",__FILE__,__LINE__); fprintf(stderr,__VA_ARGS__); fputc('\n',stderr);} }while(0) + +int main(void){ + coli_fp8_lut_init(); + + /* exhaustive round-trip on every code except the NaN patterns (0x7F/0xFF, + * decoded as 0 by design, never produced by the encoder) */ + for(int b=0;b<256;b++){ + if((b&0x7F)==0x7F) continue; /* NaN codes: decoded as 0, never produced */ + uint8_t e=coli_fp8_enc(coli_fp8_lut[b]); + CHECK(e==(uint8_t)b, "roundtrip byte 0x%02x -> %g -> 0x%02x", b, coli_fp8_lut[b], e); + } + + /* saturation and specials */ + CHECK(coli_fp8_enc(448.f)==0x7E, "448 is the max normal"); + CHECK(coli_fp8_enc(-448.f)==0xFE, "-448"); + CHECK(coli_fp8_enc(1000.f)==0x7E, "overflow saturates"); + CHECK(coli_fp8_enc(-1e30f)==0xFE, "negative overflow saturates"); + CHECK(coli_fp8_enc(INFINITY)==0x7E, "+inf saturates"); + CHECK(coli_fp8_enc(NAN)==0x00, "NaN stored as inert 0"); + CHECK(coli_fp8_enc(0.f)==0x00 && coli_fp8_enc(-0.f)==0x80, "signed zeros preserved"); + + /* round-to-nearest-even at midpoints: steps of 2 in [16,32) (E=11). + * 17 ties between 16 (mant 0, even) and 18 (mant 1, odd) -> 16; + * 19 ties between 18 (mant 1) and 20 (mant 2, even) -> 20. */ + CHECK(coli_fp8_lut[coli_fp8_enc(17.f)]==16.f, "RNE tie 17 -> 16, got %g", coli_fp8_lut[coli_fp8_enc(17.f)]); + CHECK(coli_fp8_lut[coli_fp8_enc(19.f)]==20.f, "RNE tie 19 -> 20, got %g", coli_fp8_lut[coli_fp8_enc(19.f)]); + CHECK(coli_fp8_lut[coli_fp8_enc(17.1f)]==18.f, "17.1 -> 18"); + /* denormal boundary: half of the smallest denormal (2^-10) ties to 0 (even) */ + CHECK(coli_fp8_enc(0x1p-10f)==0x00, "2^-10 ties to even 0"); + CHECK(coli_fp8_enc(0x1.8p-10f)==0x01, "0.75*2^-9 rounds up to the first denormal"); + /* mantissa overflow across the exponent boundary: just below 2^-6 */ + CHECK(coli_fp8_lut[coli_fp8_enc(0x1.fcp-7f)]==0x1p-6f, "denormal grid rounds up into the first normal"); + /* binade-boundary tie WITH mantissa carry (k==16 -> k=8,E++): 15.5 lies exactly + * between 15 (M=7, odd) and 16 (next binade, M=0, even) -> RNE must carry to 16. + * The exhaustive roundtrip never triggers this (grid points encode carry-free). */ + CHECK(coli_fp8_lut[coli_fp8_enc(15.5f)]==16.f, "RNE tie 15.5 -> 16 (mantissa carry), got %g", + coli_fp8_lut[coli_fp8_enc(15.5f)]); + CHECK(coli_fp8_lut[coli_fp8_enc(15.9f)]==16.f, "15.9 rounds up across the binade"); + CHECK(coli_fp8_lut[coli_fp8_enc(0x1.dp3f)]==14.f, "14.5 ties down to even 14"); + + /* per-row quantizer: amax maps exactly to ±448*scale, error bounded by + * |x|/16 (3 mantissa bits) + half a denormal step (scale*2^-10) */ + enum { N=576 }; + float row[N]; uint8_t q[N]; float deq[N]; + for(int i=0;i bound %g", i, deq[i], row[i], err, bound); + } + + /* all-zero row: bytes 0, scale 1 (no fabricated 1/0) */ + float zrow[8]={0}; uint8_t zq[8]; + CHECK(coli_kv8_quant_row(zrow,zq,8)==1.f, "zero row keeps scale 1"); + for(int i=0;i<8;i++) CHECK(zq[i]==0, "zero row byte %d", i); + /* subnormal amax: 448/amax would overflow to +inf and saturate every nonzero + * element to ±448 — the guard must treat the row as zero instead */ + float tiny[4]={1e-38f,-1e-38f,0,1e-40f}; uint8_t tq[4]; + CHECK(coli_kv8_quant_row(tiny,tq,4)==1.f, "subnormal row keeps scale 1"); + for(int i=0;i<4;i++) CHECK(tq[i]==0, "subnormal row byte %d", i); + + /* row accessor twin: same arithmetic as coli_kv_row, byte-typed */ + uint8_t buf[7*5]; + CHECK(coli_kv_row8(buf,4,7)==buf+28, "coli_kv_row8 arithmetic"); + CHECK(coli_kv_row8(buf,0,7)==buf, "coli_kv_row8 row 0"); + + if(fails){ fprintf(stderr,"%d failure(s)\n",fails); return 1; } + printf("OK kv_fp8 e4m3 (exhaustive roundtrip + RNE + row quant)\n"); + return 0; +} diff --git a/docs/ENVIRONMENT.md b/docs/ENVIRONMENT.md index dab656bb0..17cade861 100644 --- a/docs/ENVIRONMENT.md +++ b/docs/ENVIRONMENT.md @@ -26,6 +26,7 @@ Format: `VAR` — default — effect. | `SEED` | unset → seeded from clock + PID | RNG seed for sampling. **Unset = different every run.** Set a fixed value for reproducible sampling. | | `KVSAVE` | `1` (on) | Persist the KV cache to `/.coli_kv` so a conversation reopens warm. `KVSAVE=0` disables save+load (lossless round-trip; does not change output). | | `KV_SLOTS` | `1` | Number of independent KV conversation slots (1–16), used in serve mode. | +| `KV8` | `0` (off) | Store the MLA latent KV cache in fp8 e4m3 with a per-row scale: ~3.9× less KV RAM, and `.coli_kv` shrinks ~4× (saved as the v2 format; f32 v1 files are quantized on resume and rewritten). Adds DeepSeek-V3-class KV quantization noise to attention. Covered on the CPU and CUDA attention paths (fp8 absorb kernels, ¼ the PCIe traffic); auto-disabled under `COLI_METAL`, and forces `COLI_CUDA_PIPE=0`. | | `THINK` | `0` (off) | Emit a `` reasoning block. `THINK=1` turns on visible reasoning. | | `MTP` | on | Multi-Token Prediction (speculative draft head). `MTP=0` disables it. |