From 6826c00a6dc286780367af3824259d747efad6af Mon Sep 17 00:00:00 2001 From: monotophic Date: Thu, 3 Sep 2026 12:56:54 -0400 Subject: [PATCH 1/6] refactor(quant): move FP8_BLOCK/fp8_nblk into fp8_format.h The CUDA backend needs the same 128-column block edge and block count the CPU decoder uses. Sharing them through quant.h is not possible: that header pulls in the whole CPU kernel set, which backend_cuda.cu must not compile. Move the two definitions into a minimal header both sides include, so there is one definition of the fp8 block geometry rather than a constant repeated per backend. Every prerequisite list whose translation unit reaches the new header gains it beside quant.h, in both build files. One rule, tests/test_qwen36_dnproj_batch, names quant.h but is left alone: qwen36.c does not include quant.h, so that translation unit never reaches fp8_format.h and listing it would declare a dependency that does not exist. That claim is now enforced rather than asserted. The existing Makefile prerequisite test walks only an engine's own #include lines, so a header reached THROUGH another one is invisible to it -- which is exactly fp8_format.h's shape, since quant.h includes it and no engine .c does. A new case closes the include graph for this header: drop fp8_format.h from a rule whose engine reaches it and the test names that rule. It is scoped to this header deliberately; the same closure over every header reports seven rules that predate this change, which the case documents as a separate fix rather than failing on arrival for reasons this change did not cause. Pointers to the old location (the fmt=8 entry in docs/FORMATS.md, comments in colibri.c and backend_metal.mm) now name fp8_format.h. FORMATS.md's source references drop their line numbers and cite symbols instead: those anchors went stale twice, and a stale line number reads as a verification that was not performed. --- c/Makefile | 128 +++++++++++++++++----------------- c/Makefile.deepseek-v4 | 6 +- c/backend_metal.mm | 2 +- c/colibri.c | 7 +- c/fp8_format.h | 30 ++++++++ c/quant.h | 6 +- c/tests/test_makefile_deps.py | 52 ++++++++++++++ docs/FORMATS.md | 58 +++++++-------- 8 files changed, 187 insertions(+), 102 deletions(-) create mode 100644 c/fp8_format.h diff --git a/c/Makefile b/c/Makefile index 4842b9fdf..9fedfa799 100644 --- a/c/Makefile +++ b/c/Makefile @@ -765,7 +765,7 @@ $(file >.build-config,$(BUILD_CONFIG)) endif .build-config: ; -colibri$(EXE): colibri.c oracle.h pin_pool.h cli_args.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h quant.h idot.h sample.h kv_persist.h telemetry.h route_trace.h omp_tune.h kv_fp8.h kv_tq.h abl.h backend_cuda.h backend_metal.h backend_vulkan.h decode_batch.h edge_adapters.h edge_runtime.h edge_tok_internal.h schema_gbnf.h segment_adapter_internal.h segment_adapters.h segment_runtime.h tier.h $(CUDA_OBJ) $(METAL_OBJ) $(VK_OBJ) $(VK_SPV) .build-config +colibri$(EXE): colibri.c oracle.h pin_pool.h cli_args.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h route_trace.h omp_tune.h kv_fp8.h kv_tq.h abl.h backend_cuda.h backend_metal.h backend_vulkan.h decode_batch.h edge_adapters.h edge_runtime.h edge_tok_internal.h schema_gbnf.h segment_adapter_internal.h segment_adapters.h segment_runtime.h tier.h $(CUDA_OBJ) $(METAL_OBJ) $(VK_OBJ) $(VK_SPV) .build-config $(CC) $(CFLAGS) colibri.c $(CUDA_OBJ) $(METAL_OBJ) $(VK_OBJ) -o colibri$(EXE) $(LDFLAGS) # Vulkan backend object (plain C + vulkan headers) and its SPIR-V shaders. @@ -1123,7 +1123,7 @@ qwen36$(EXE): qwen36.c decode_batch.h serve_poll.h cli_args.h qwen36_tier.h expe # DeepSeek V4.1 Flash: one file, like every other portable engine. The fp4 experts # stream from the official checkpoint and quant.h's mxfp4 kernel reads them as they # are, so there is no conversion target here to go with it. -deepseek_v41$(EXE): deepseek_v41.c cli_args.h st.h json.h tok.h tok_unicode.h compat.h omp_tune.h kv_prefix.h pin_pool.h quant.h idot.h \ +deepseek_v41$(EXE): deepseek_v41.c cli_args.h st.h json.h tok.h tok_unicode.h compat.h omp_tune.h kv_prefix.h pin_pool.h quant.h fp8_format.h idot.h \ sparse_attn.h hyper_connections.h serve_codec.h serve_poll.h .build-config $(CC) $(CFLAGS) deepseek_v41.c -o deepseek_v41$(EXE) $(LDFLAGS) @@ -1131,7 +1131,7 @@ deepseek_v41$(EXE): deepseek_v41.c cli_args.h st.h json.h tok.h tok_unicode.h co # loader and SERVE=1 protocol. With CUDA=1 it links the same expert tier as # qwen36 (fp8 streaming mode: hot experts get VRAM copies, the RAM LRU stays); # without it the tier header's inline stubs keep the build toolkit-free. -qwen38$(EXE): qwen38.c pin_pool.h cli_args.h qwen38_core.h kv_prefix.h qwen38_nfc.h qwen38_nfc_tables.h st.h json.h compat.h omp_tune.h quant.h idot.h route_trace.h tok.h tok_unicode.h tok_unicode_o200k.h serve_codec.h edge_adapter_internal.h edge_adapters.h edge_runtime.h qwen38_vision.h segment_adapter_internal.h segment_adapters.h segment_runtime.h qwen36_tier.h $(QWEN36_TIER_SRC) $(CUDA_OBJ) +qwen38$(EXE): qwen38.c pin_pool.h cli_args.h qwen38_core.h kv_prefix.h qwen38_nfc.h qwen38_nfc_tables.h st.h json.h compat.h omp_tune.h quant.h fp8_format.h idot.h route_trace.h tok.h tok_unicode.h tok_unicode_o200k.h serve_codec.h edge_adapter_internal.h edge_adapters.h edge_runtime.h qwen38_vision.h segment_adapter_internal.h segment_adapters.h segment_runtime.h qwen36_tier.h $(QWEN36_TIER_SRC) $(CUDA_OBJ) $(CC) $(QWEN36_CFLAGS) qwen38.c $(QWEN36_TIER_SRC) $(CUDA_OBJ) -o qwen38$(EXE) $(QWEN36_LDFLAGS) .PHONY: qwen38-tiny-generate qwen38-tiny-check @@ -1242,10 +1242,10 @@ NOCUDA_LDFLAGS = $(filter-out -lcudart -lstdc++ -lcuda -L$(CUDA_HOME)/lib64 \ # GLM-5.3-Flash: routed experts stream from the int4-gs64 container. # METAL=1 accelerates resident matrices and routed MoE; CPU remains fallback. -glm53$(EXE): glm53.c decode_batch.h pin_pool.h cli_args.h st.h json.h stop_ids.h tok.h tok_unicode.h tok_unicode_o200k.h compat.h omp_tune.h serve_poll.h route_trace.h quant.h idot.h hyper_connections.h delta_attention.h sparse_index.h vision_tower.h backend_metal.h backend_vulkan.h edge_adapter_internal.h edge_adapters.h edge_runtime.h segment_adapter_internal.h segment_adapters.h segment_runtime.h $(METAL_OBJ) $(VK_OBJ) $(VK_SPV) +glm53$(EXE): glm53.c decode_batch.h pin_pool.h cli_args.h st.h json.h stop_ids.h tok.h tok_unicode.h tok_unicode_o200k.h compat.h omp_tune.h serve_poll.h route_trace.h quant.h fp8_format.h idot.h hyper_connections.h delta_attention.h sparse_index.h vision_tower.h backend_metal.h backend_vulkan.h edge_adapter_internal.h edge_adapters.h edge_runtime.h segment_adapter_internal.h segment_adapters.h segment_runtime.h $(METAL_OBJ) $(VK_OBJ) $(VK_SPV) $(CC) $(CFLAGS) glm53.c $(METAL_OBJ) $(VK_OBJ) -o glm53$(EXE) $(LDFLAGS) -kimi_k3$(EXE): kimi_k3.c cli_args.h st.h json.h tok.h tok_unicode.h tok_unicode_o200k.h compat.h quant.h idot.h omp_tune.h route_trace.h kv_prefix.h pin_pool.h serve_codec.h backend_cuda.h backend_metal.h backend_vulkan.h edge_adapters.h edge_runtime.h edge_tok_internal.h hybrid_split.h segment_adapter_internal.h segment_adapters.h segment_runtime.h $(CUDA_OBJ) $(VK_OBJ) $(VK_SPV) $(METAL_OBJ) +kimi_k3$(EXE): kimi_k3.c cli_args.h st.h json.h tok.h tok_unicode.h tok_unicode_o200k.h compat.h quant.h fp8_format.h idot.h omp_tune.h route_trace.h kv_prefix.h pin_pool.h serve_codec.h backend_cuda.h backend_metal.h backend_vulkan.h edge_adapters.h edge_runtime.h edge_tok_internal.h hybrid_split.h segment_adapter_internal.h segment_adapters.h segment_runtime.h $(CUDA_OBJ) $(VK_OBJ) $(VK_SPV) $(METAL_OBJ) $(CC) $(CFLAGS) kimi_k3.c $(CUDA_OBJ) $(VK_OBJ) $(METAL_OBJ) -o kimi_k3$(EXE) $(LDFLAGS) # Use a baseline that matches the compiler target. macOS already targets a @@ -1273,7 +1273,7 @@ iobench$(EXE): iobench.c compat.h tests/test_serve_sentinel$(EXE): tests/test_serve_sentinel.c compat.h serve_codec.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) -tests/test_cluster_protocol$(EXE): tests/test_cluster_protocol.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h idot.h sample.h kv_persist.h telemetry.h route_trace.h +tests/test_cluster_protocol$(EXE): tests/test_cluster_protocol.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h route_trace.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) tests/test_ue8m0$(EXE): tests/test_ue8m0.c st.h json.h compat.h @@ -1288,10 +1288,10 @@ tests/test_tok_o200k$(EXE): tests/test_tok_o200k.c tok.h tok_unicode.h tok_unico tests/test_tok_gpt2$(EXE): tests/test_tok_gpt2.c tok.h tok_unicode.h tok_unicode_o200k.h json.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) -tests/test_k3_ram_budget$(EXE): tests/test_k3_ram_budget.c kimi_k3.c st.h tok.h quant.h idot.h compat.h omp_tune.h route_trace.h kv_prefix.h serve_codec.h +tests/test_k3_ram_budget$(EXE): tests/test_k3_ram_budget.c kimi_k3.c st.h tok.h quant.h fp8_format.h idot.h compat.h omp_tune.h route_trace.h kv_prefix.h serve_codec.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) -tests/test_k3_mmap$(EXE): tests/test_k3_mmap.c kimi_k3.c st.h tok.h quant.h idot.h compat.h omp_tune.h route_trace.h kv_prefix.h serve_codec.h +tests/test_k3_mmap$(EXE): tests/test_k3_mmap.c kimi_k3.c st.h tok.h quant.h fp8_format.h idot.h compat.h omp_tune.h route_trace.h kv_prefix.h serve_codec.h $(CC) $(NOCUDA_CFLAGS) $< $(VK_OBJ) -o $@ $(NOCUDA_LDFLAGS) tests/test_tok_kimi_tiny$(EXE): tests/test_tok_kimi_tiny.c tok.h tok_unicode.h tok_unicode_o200k.h json.h @@ -1302,19 +1302,19 @@ tests/test_st_pread$(EXE): tests/test_st_pread.c st.h json.h compat.h tests/test_st_slice$(EXE): tests/test_st_slice.c st.h json.h compat.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) -tests/test_qwen38_tokenizer$(EXE): tests/test_qwen38_tokenizer.c qwen38.c qwen38_core.h kv_prefix.h qwen38_nfc.h qwen38_nfc_tables.h edge_runtime.c edge_runtime.h edge_adapters.h edge_adapter_internal.h st.h json.h compat.h quant.h idot.h route_trace.h tok_unicode.h tok_unicode_o200k.h +tests/test_qwen38_tokenizer$(EXE): tests/test_qwen38_tokenizer.c qwen38.c qwen38_core.h kv_prefix.h qwen38_nfc.h qwen38_nfc_tables.h edge_runtime.c edge_runtime.h edge_adapters.h edge_adapter_internal.h st.h json.h compat.h quant.h fp8_format.h idot.h route_trace.h tok_unicode.h tok_unicode_o200k.h $(CC) $(NOCUDA_CFLAGS) $< edge_runtime.c -o $@ $(NOCUDA_LDFLAGS) -tests/test_qwen38_config$(EXE): tests/test_qwen38_config.c qwen38.c qwen38_core.h kv_prefix.h qwen38_nfc.h qwen38_nfc_tables.h st.h json.h compat.h quant.h idot.h route_trace.h tok_unicode.h tok_unicode_o200k.h +tests/test_qwen38_config$(EXE): tests/test_qwen38_config.c qwen38.c qwen38_core.h kv_prefix.h qwen38_nfc.h qwen38_nfc_tables.h st.h json.h compat.h quant.h fp8_format.h idot.h route_trace.h tok_unicode.h tok_unicode_o200k.h $(CC) $(NOCUDA_CFLAGS) $< -o $@ $(NOCUDA_LDFLAGS) -tests/test_dsv41_serve_budget$(EXE): tests/test_dsv41_serve_budget.c deepseek_v41.c cli_args.h st.h json.h tok.h tok_unicode.h compat.h omp_tune.h kv_prefix.h pin_pool.h quant.h idot.h sparse_attn.h hyper_connections.h serve_codec.h serve_poll.h +tests/test_dsv41_serve_budget$(EXE): tests/test_dsv41_serve_budget.c deepseek_v41.c cli_args.h st.h json.h tok.h tok_unicode.h compat.h omp_tune.h kv_prefix.h pin_pool.h quant.h fp8_format.h idot.h sparse_attn.h hyper_connections.h serve_codec.h serve_poll.h $(CC) $(NOCUDA_CFLAGS) $< -o $@ $(NOCUDA_LDFLAGS) -tests/test_qwen38_idot$(EXE): tests/test_qwen38_idot.c qwen38.c qwen38_core.h kv_prefix.h qwen38_nfc.h qwen38_nfc_tables.h st.h json.h compat.h quant.h idot.h route_trace.h tok_unicode.h tok_unicode_o200k.h +tests/test_qwen38_idot$(EXE): tests/test_qwen38_idot.c qwen38.c qwen38_core.h kv_prefix.h qwen38_nfc.h qwen38_nfc_tables.h st.h json.h compat.h quant.h fp8_format.h idot.h route_trace.h tok_unicode.h tok_unicode_o200k.h $(CC) $(NOCUDA_CFLAGS) $< -o $@ $(NOCUDA_LDFLAGS) -tests/test_qwen38_serve_framing$(EXE): tests/test_qwen38_serve_framing.c qwen38.c qwen38_core.h kv_prefix.h qwen38_nfc.h qwen38_nfc_tables.h st.h json.h compat.h quant.h idot.h route_trace.h serve_codec.h tok_unicode.h tok_unicode_o200k.h +tests/test_qwen38_serve_framing$(EXE): tests/test_qwen38_serve_framing.c qwen38.c qwen38_core.h kv_prefix.h qwen38_nfc.h qwen38_nfc_tables.h st.h json.h compat.h quant.h fp8_format.h idot.h route_trace.h serve_codec.h tok_unicode.h tok_unicode_o200k.h $(CC) $(NOCUDA_CFLAGS) $< -o $@ $(NOCUDA_LDFLAGS) tests/test_qwen38_vision$(EXE): tests/test_qwen38_vision.c qwen38_vision.h st.h json.h compat.h @@ -1336,13 +1336,13 @@ qwen38-vision-serve-check: qwen38$(EXE) $(PYTHON) tools/make_edge_tiny_tokenizer.py --vocab-size 64 ./qwen38_mm_tiny $(PYTHON) tests/test_qwen38_vision_serve.py --binary ./qwen38$(EXE) --fixture ./qwen38_mm_tiny -tests/test_qwen38_prefix$(EXE): tests/test_qwen38_prefix.c qwen38.c qwen38_core.h kv_prefix.h qwen38_nfc.h qwen38_nfc_tables.h st.h json.h compat.h quant.h idot.h route_trace.h serve_codec.h tok_unicode.h tok_unicode_o200k.h +tests/test_qwen38_prefix$(EXE): tests/test_qwen38_prefix.c qwen38.c qwen38_core.h kv_prefix.h qwen38_nfc.h qwen38_nfc_tables.h st.h json.h compat.h quant.h fp8_format.h idot.h route_trace.h serve_codec.h tok_unicode.h tok_unicode_o200k.h $(CC) $(NOCUDA_CFLAGS) $< -o $@ $(NOCUDA_LDFLAGS) -tests/test_qwen38_metrics$(EXE): tests/test_qwen38_metrics.c qwen38.c qwen38_core.h kv_prefix.h qwen38_nfc.h qwen38_nfc_tables.h st.h json.h compat.h quant.h idot.h route_trace.h serve_codec.h tok_unicode.h tok_unicode_o200k.h +tests/test_qwen38_metrics$(EXE): tests/test_qwen38_metrics.c qwen38.c qwen38_core.h kv_prefix.h qwen38_nfc.h qwen38_nfc_tables.h st.h json.h compat.h quant.h fp8_format.h idot.h route_trace.h serve_codec.h tok_unicode.h tok_unicode_o200k.h $(CC) $(NOCUDA_CFLAGS) $< -o $@ $(NOCUDA_LDFLAGS) -tests/test_qwen38_native_weights$(EXE): tests/test_qwen38_native_weights.c qwen38.c qwen38_core.h kv_prefix.h qwen38_nfc.h qwen38_nfc_tables.h segment_runtime.c segment_runtime.h segment_adapters.h segment_adapter_internal.h st.h json.h compat.h quant.h idot.h route_trace.h serve_codec.h tok_unicode.h tok_unicode_o200k.h +tests/test_qwen38_native_weights$(EXE): tests/test_qwen38_native_weights.c qwen38.c qwen38_core.h kv_prefix.h qwen38_nfc.h qwen38_nfc_tables.h segment_runtime.c segment_runtime.h segment_adapters.h segment_adapter_internal.h st.h json.h compat.h quant.h fp8_format.h idot.h route_trace.h serve_codec.h tok_unicode.h tok_unicode_o200k.h $(CC) $(NOCUDA_CFLAGS) $< segment_runtime.c -o $@ $(NOCUDA_LDFLAGS) tests/test_st_map$(EXE): tests/test_st_map.c st.h json.h compat.h @@ -1388,12 +1388,12 @@ tests/test_grammar$(EXE): tests/test_grammar.c grammar.h # equivalence under the shipping flags is covered by the differential dump # (old vs new rope_interleave are byte-identical); this test guards the # source-level formula equivalence and must not depend on contraction luck. -tests/test_rope_invfreq$(EXE): tests/test_rope_invfreq.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h idot.h sample.h kv_persist.h telemetry.h route_trace.h +tests/test_rope_invfreq$(EXE): tests/test_rope_invfreq.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h route_trace.h $(CC) $(CFLAGS) -ffp-contract=off $< -o $@ $(LDFLAGS) # schema->GBNF compile cache (#7): grammar_reset must equal a fresh setup, and # the GrDraft.src ownership must not leak/double-free. Includes colibri.c. -tests/test_grammar_cache$(EXE): tests/test_grammar_cache.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h idot.h sample.h kv_persist.h telemetry.h route_trace.h +tests/test_grammar_cache$(EXE): tests/test_grammar_cache.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h route_trace.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) # Standalone: drives a faithful miniature of moe()'s routing+accumulate and links @@ -1409,7 +1409,7 @@ tests/test_degrade_zero$(EXE): tests/test_degrade_zero.c tests/test_schema_gbnf$(EXE): tests/test_schema_gbnf.c schema_gbnf.h grammar.h json.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) -tests/test_spec_decode_state$(EXE): tests/test_spec_decode_state.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h idot.h sample.h kv_persist.h telemetry.h route_trace.h +tests/test_spec_decode_state$(EXE): tests/test_spec_decode_state.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h route_trace.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) tests/test_decode_batch$(EXE): tests/test_decode_batch.c decode_batch.h @@ -1441,13 +1441,13 @@ tests/bench_inkling_shared_batch$(EXE): tests/bench_inkling_shared_batch.c inkli tests/test_inkling_cache_index$(EXE): tests/test_inkling_cache_index.c inkling.c st.h json.h tok.h tok_unicode.h tok_unicode_o200k.h compat.h omp_tune.h route_trace.h kv_prefix.h serve_codec.h $(INK_CUDA_OBJ) $(METAL_OBJ) $(CC) $(CFLAGS) $< $(INK_CUDA_OBJ) $(METAL_OBJ) -o $@ $(LDFLAGS) -tests/test_kimi_serve_framing$(EXE): tests/test_kimi_serve_framing.c kimi_k3.c kv_prefix.h serve_codec.h st.h json.h tok.h tok_unicode.h tok_unicode_o200k.h compat.h quant.h idot.h omp_tune.h route_trace.h +tests/test_kimi_serve_framing$(EXE): tests/test_kimi_serve_framing.c kimi_k3.c kv_prefix.h serve_codec.h st.h json.h tok.h tok_unicode.h tok_unicode_o200k.h compat.h quant.h fp8_format.h idot.h omp_tune.h route_trace.h $(CC) $(NOCUDA_CFLAGS) $< $(VK_OBJ) -o $@ $(NOCUDA_LDFLAGS) -tests/test_kimi_cache_index$(EXE): tests/test_kimi_cache_index.c kimi_k3.c kv_prefix.h serve_codec.h st.h json.h tok.h tok_unicode.h tok_unicode_o200k.h compat.h quant.h idot.h omp_tune.h route_trace.h +tests/test_kimi_cache_index$(EXE): tests/test_kimi_cache_index.c kimi_k3.c kv_prefix.h serve_codec.h st.h json.h tok.h tok_unicode.h tok_unicode_o200k.h compat.h quant.h fp8_format.h idot.h omp_tune.h route_trace.h $(CC) $(NOCUDA_CFLAGS) $< $(VK_OBJ) -o $@ $(NOCUDA_LDFLAGS) -tests/test_k3_chat_tools$(EXE): tests/test_k3_chat_tools.c kimi_k3.c kv_prefix.h serve_codec.h st.h json.h tok.h tok_unicode.h tok_unicode_o200k.h compat.h quant.h idot.h omp_tune.h route_trace.h +tests/test_k3_chat_tools$(EXE): tests/test_k3_chat_tools.c kimi_k3.c kv_prefix.h serve_codec.h st.h json.h tok.h tok_unicode.h tok_unicode_o200k.h compat.h quant.h fp8_format.h idot.h omp_tune.h route_trace.h $(CC) $(NOCUDA_CFLAGS) $< $(VK_OBJ) -o $@ $(NOCUDA_LDFLAGS) # olmoe's matmul_q, not colibri's: compares the IDOT path against FP32 @@ -1465,33 +1465,33 @@ tests/test_olmoe_cache_index$(EXE): tests/test_olmoe_cache_index.c olmoe.c st.h tests/test_qwen36_cache_index$(EXE): tests/test_qwen36_cache_index.c qwen36.c expert_ffn.h st.h json.h compat.h sample.h tok.h tok_unicode.h tok_unicode_o200k.h omp_tune.h route_trace.h serve_codec.h qwen36_tier.h $(QWEN36_TIER_SRC) $(CUDA_OBJ) $(CC) $(QWEN36_CFLAGS) $< $(QWEN36_TIER_SRC) $(CUDA_OBJ) -o $@ $(QWEN36_LDFLAGS) -tests/test_idot$(EXE): tests/test_idot.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h idot.h sample.h kv_persist.h telemetry.h route_trace.h +tests/test_idot$(EXE): tests/test_idot.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h route_trace.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) -tests/test_i4_grouped$(EXE): tests/test_i4_grouped.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h idot.h sample.h kv_persist.h telemetry.h route_trace.h +tests/test_i4_grouped$(EXE): tests/test_i4_grouped.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h route_trace.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) -tests/test_stops$(EXE): tests/test_stops.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h idot.h sample.h kv_persist.h telemetry.h route_trace.h +tests/test_stops$(EXE): tests/test_stops.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h route_trace.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) -tests/test_cfg_topk$(EXE): tests/test_cfg_topk.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h idot.h sample.h kv_persist.h telemetry.h route_trace.h +tests/test_cfg_topk$(EXE): tests/test_cfg_topk.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h route_trace.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) -tests/test_topp$(EXE): tests/test_topp.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h idot.h sample.h kv_persist.h telemetry.h route_trace.h +tests/test_topp$(EXE): tests/test_topp.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h route_trace.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) # bench_topp is a microbenchmark (old qsort vs new heap partial-select, #335), NOT a test # gate -- intentionally absent from TEST_BINS. Build on demand: make tests/bench_topp -tests/bench_topp$(EXE): tests/bench_topp.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h idot.h sample.h kv_persist.h telemetry.h route_trace.h +tests/bench_topp$(EXE): tests/bench_topp.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h route_trace.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) -tests/test_sample_nan$(EXE): tests/test_sample_nan.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h idot.h sample.h kv_persist.h telemetry.h route_trace.h +tests/test_sample_nan$(EXE): tests/test_sample_nan.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h route_trace.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) -tests/test_temp_env$(EXE): tests/test_temp_env.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h idot.h sample.h kv_persist.h telemetry.h route_trace.h +tests/test_temp_env$(EXE): tests/test_temp_env.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h route_trace.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) -tests/test_kv_alloc$(EXE): tests/test_kv_alloc.c colibri.c oracle.h st.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h idot.h sample.h kv_persist.h telemetry.h route_trace.h kv_fp8.h kv_tq.h +tests/test_kv_alloc$(EXE): tests/test_kv_alloc.c colibri.c oracle.h st.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h route_trace.h kv_fp8.h kv_tq.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) tests/test_kv_fp8$(EXE): tests/test_kv_fp8.c kv_fp8.h decode_batch.h @@ -1500,15 +1500,15 @@ tests/test_kv_fp8$(EXE): tests/test_kv_fp8.c kv_fp8.h decode_batch.h tests/test_kv_tq$(EXE): tests/test_kv_tq.c kv_tq.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) -tests/test_kv_disk$(EXE): tests/test_kv_disk.c colibri.c oracle.h st.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h idot.h sample.h kv_persist.h telemetry.h route_trace.h kv_fp8.h kv_tq.h +tests/test_kv_disk$(EXE): tests/test_kv_disk.c colibri.c oracle.h st.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h route_trace.h kv_fp8.h kv_tq.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) # fmt=6 kernel oracle: needs the generated grid table, and its fixture comes from # the reference codec (tools/make_e8_fixture.py) — regenerate if the layout moves. -tests/test_e8_kernel$(EXE): tests/test_e8_kernel.c quant.h idot.h +tests/test_e8_kernel$(EXE): tests/test_e8_kernel.c quant.h fp8_format.h idot.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) -tests/test_e4m3_vector$(EXE): tests/test_e4m3_vector.c quant.h idot.h +tests/test_e4m3_vector$(EXE): tests/test_e4m3_vector.c quant.h fp8_format.h idot.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) tests/test_stop_ids$(EXE): tests/test_stop_ids.c stop_ids.h json.h @@ -1535,13 +1535,13 @@ fuzz-rans: tests/fuzz_rans.c rans.h tests/fuzz_rans.c -o tests/fuzz_rans -lm ./tests/fuzz_rans -tests/test_int3$(EXE): tests/test_int3.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h idot.h sample.h kv_persist.h telemetry.h route_trace.h +tests/test_int3$(EXE): tests/test_int3.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h route_trace.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) -tests/test_int3_load$(EXE): tests/test_int3_load.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h idot.h sample.h kv_persist.h telemetry.h route_trace.h +tests/test_int3_load$(EXE): tests/test_int3_load.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h route_trace.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) -tests/test_fp8_passthrough$(EXE): tests/test_fp8_passthrough.c quant.h idot.h +tests/test_fp8_passthrough$(EXE): tests/test_fp8_passthrough.c quant.h fp8_format.h idot.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) # Host-only: backend_cuda.h's format predicate is plain C, so its truth table is @@ -1550,10 +1550,10 @@ tests/test_fp8_passthrough$(EXE): tests/test_fp8_passthrough.c quant.h idot.h tests/test_cuda_fmt_guard$(EXE): tests/test_cuda_fmt_guard.c backend_cuda.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) -tests/test_fp8_load$(EXE): tests/test_fp8_load.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h idot.h sample.h kv_persist.h telemetry.h +tests/test_fp8_load$(EXE): tests/test_fp8_load.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) -tests/test_qt_addrow$(EXE): tests/test_qt_addrow.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h idot.h sample.h kv_persist.h telemetry.h +tests/test_qt_addrow$(EXE): tests/test_qt_addrow.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) tests/test_logit_nan$(EXE): tests/test_logit_nan.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h @@ -1571,7 +1571,7 @@ tests/test_compat_direct$(EXE): tests/test_compat_direct.c compat.h tests/test_expert_store_ops$(EXE): tests/test_expert_store_ops.c expert_store.h tensor.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) -tests/test_native_quant$(EXE): tests/test_native_quant.c deepseek_v4.c native_quant.h tensor.h quant.h idot.h +tests/test_native_quant$(EXE): tests/test_native_quant.c deepseek_v4.c native_quant.h tensor.h quant.h fp8_format.h idot.h $(CC) $(CFLAGS) -DCOLI_V4_UNIT_NATIVE_QUANT deepseek_v4.c $< -o $@ $(LDFLAGS) tests/test_edge_runtime$(EXE): tests/test_edge_runtime.c edge_runtime.c edge_runtime.h @@ -1616,7 +1616,7 @@ $(SEGMENT_BUILD_DIR)/glm.o: colibri.c oracle.h segment_runtime.h edge_runtime.h $(SEGMENT_BUILD_DIR)/glm53.o: glm53.c segment_runtime.h edge_runtime.h \ segment_adapters.h edge_adapters.h segment_adapter_internal.h \ - edge_adapter_internal.h st.h quant.h tok.h hyper_connections.h \ + edge_adapter_internal.h st.h quant.h fp8_format.h tok.h hyper_connections.h \ delta_attention.h sparse_index.h vision_tower.h | $(SEGMENT_BUILD_DIR) $(CC) $(SEGMENT_CPU_CFLAGS) -DGLM53_NO_MAIN -c glm53.c -o $@ @@ -1642,7 +1642,7 @@ $(SEGMENT_BUILD_DIR)/qwen36.o: qwen36.c segment_runtime.h edge_runtime.h \ $(SEGMENT_BUILD_DIR)/qwen38.o: qwen38.c qwen38_core.h kv_prefix.h segment_runtime.h \ segment_adapters.h segment_adapter_internal.h edge_runtime.h \ - edge_adapters.h edge_adapter_internal.h st.h json.h compat.h quant.h \ + edge_adapters.h edge_adapter_internal.h st.h json.h compat.h quant.h fp8_format.h \ qwen38_nfc.h qwen38_nfc_tables.h route_trace.h tok_unicode.h \ tok_unicode_o200k.h | $(SEGMENT_BUILD_DIR) $(CC) $(SEGMENT_CPU_CFLAGS) -DQWEN38_NO_MAIN -c qwen38.c -o $@ @@ -1808,7 +1808,7 @@ $(V4_OWN_DIR)/COLI_V4_UNIT_CONFIG.o: deepseek_v4.c deepseek_v4.h deepseek_v4_int $(V4_OWN_DIR)/COLI_V4_UNIT_ST.o: deepseek_v4.c deepseek_v4.h deepseek_v4_internal.h st.h | $(V4_OWN_DIR) $(CC) $(V4_OWN_CFLAGS) -DCOLI_V4_UNIT_ST -c deepseek_v4.c -o $@ -$(V4_OWN_DIR)/COLI_V4_UNIT_NATIVE_QUANT.o: deepseek_v4.c deepseek_v4.h deepseek_v4_internal.h quant.h idot.h | $(V4_OWN_DIR) +$(V4_OWN_DIR)/COLI_V4_UNIT_NATIVE_QUANT.o: deepseek_v4.c deepseek_v4.h deepseek_v4_internal.h quant.h fp8_format.h idot.h | $(V4_OWN_DIR) $(CC) $(V4_OWN_CFLAGS) -DCOLI_V4_UNIT_NATIVE_QUANT -c deepseek_v4.c -o $@ # The engine (RUNTIME unit) opens its expert store through the pluggable @@ -1937,7 +1937,7 @@ tests/test_qwen36_tier_int8_decode$(EXE): tests/test_qwen36_tier_int8_decode.c t # The qwen38 engine through its own main() on the fake backend: tier start # from the FP8 fixture, reduced CPU list, qt_note on recycled slots, oracle # tokens unchanged. Needs the fixture (qwen38-tiny-fp8-generate). -tests/test_qwen38_tier_engine$(EXE): tests/test_qwen38_tier_engine.c tests/qwen36_fake_cuda.h qwen38.c qwen38_core.h kv_prefix.h qwen36_tier.c qwen36_tier.h backend_cuda.h cli_args.h st.h json.h compat.h quant.h idot.h route_trace.h +tests/test_qwen38_tier_engine$(EXE): tests/test_qwen38_tier_engine.c tests/qwen36_fake_cuda.h qwen38.c qwen38_core.h kv_prefix.h qwen36_tier.c qwen36_tier.h backend_cuda.h cli_args.h st.h json.h compat.h quant.h fp8_format.h idot.h route_trace.h $(CC) $(CFLAGS) -DCOLI_CUDA $< -o $@ $(LDFLAGS) tests/test_serve_poll$(EXE): tests/test_serve_poll.c serve_poll.h @@ -1956,15 +1956,15 @@ tests/test_mem_available$(EXE): tests/test_mem_available.c compat.h tests/test_798_guards$(EXE): tests/test_798_guards.c st.h json.h compat.h route_trace.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) -tests/test_dsa_select$(EXE): tests/test_dsa_select.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h idot.h sample.h kv_persist.h telemetry.h route_trace.h +tests/test_dsa_select$(EXE): tests/test_dsa_select.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h route_trace.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) -tests/test_corpus_draft$(EXE): tests/test_corpus_draft.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h idot.h sample.h kv_persist.h telemetry.h route_trace.h +tests/test_corpus_draft$(EXE): tests/test_corpus_draft.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h route_trace.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) -tests/test_cap_precedence$(EXE): tests/test_cap_precedence.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h idot.h sample.h kv_persist.h telemetry.h +tests/test_cap_precedence$(EXE): tests/test_cap_precedence.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) -tests/test_mirror_stripe_split$(EXE): tests/test_mirror_stripe_split.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h idot.h sample.h kv_persist.h telemetry.h +tests/test_mirror_stripe_split$(EXE): tests/test_mirror_stripe_split.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) tests/test_v4_hybrid_policy$(EXE): tests/test_v4_hybrid_policy.c deepseek_v4_hybrid.h hybrid_split.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) @@ -1975,29 +1975,29 @@ tests/test_k3_fill_budget$(EXE): tests/test_k3_fill_budget.c hybrid_split.h tests/test_v4_bank_pair$(EXE): tests/test_v4_bank_pair.c deepseek_v4_bank_pair.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) -tests/test_ram_clamp$(EXE): tests/test_ram_clamp.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h idot.h sample.h kv_persist.h telemetry.h +tests/test_ram_clamp$(EXE): tests/test_ram_clamp.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) -tests/test_cap_mixed_width$(EXE): tests/test_cap_mixed_width.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h idot.h sample.h kv_persist.h telemetry.h route_trace.h +tests/test_cap_mixed_width$(EXE): tests/test_cap_mixed_width.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h route_trace.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) -tests/test_eslot_inflight$(EXE): tests/test_eslot_inflight.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h idot.h sample.h kv_persist.h telemetry.h route_trace.h +tests/test_eslot_inflight$(EXE): tests/test_eslot_inflight.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h route_trace.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) -tests/test_glm_cache_index$(EXE): tests/test_glm_cache_index.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h idot.h sample.h kv_persist.h telemetry.h route_trace.h +tests/test_glm_cache_index$(EXE): tests/test_glm_cache_index.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h route_trace.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) -tests/test_ssd_probe$(EXE): tests/test_ssd_probe.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h idot.h sample.h kv_persist.h telemetry.h +tests/test_ssd_probe$(EXE): tests/test_ssd_probe.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) # bench_dsa_select is a microbenchmark (old qsort vs new quickselect partial-select, #356), # NOT a test gate -- intentionally absent from TEST_BINS. Build on demand: make tests/bench_dsa_select -tests/bench_dsa_select$(EXE): tests/bench_dsa_select.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h idot.h sample.h kv_persist.h telemetry.h route_trace.h +tests/bench_dsa_select$(EXE): tests/bench_dsa_select.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h route_trace.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) # bench_router_select is a microbenchmark (duplicate-prefix scan vs marked-score scan), # not a test gate. Build on demand: make tests/bench_router_select. -tests/bench_router_select$(EXE): tests/bench_router_select.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h idot.h sample.h kv_persist.h telemetry.h route_trace.h +tests/bench_router_select$(EXE): tests/bench_router_select.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h route_trace.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) # bench_indexer_allocations: microbenchmark (DeepSeek V4 indexer malloc vs persistent arena scratch), NOT a test gate. @@ -2007,34 +2007,34 @@ tests/bench_indexer_allocations$(EXE): tests/bench_indexer_allocations.c # bench_idot: microbenchmark (single-acc vs independent-acc AVX-VNNI idot), NOT a test gate. # Build on demand on an AVX-VNNI CPU: make tests/bench_idot ARCH=native -tests/bench_idot$(EXE): tests/bench_idot.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h idot.h sample.h kv_persist.h telemetry.h route_trace.h +tests/bench_idot$(EXE): tests/bench_idot.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h route_trace.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) # bench_i4p_gidot: microbenchmark (per-row vs multi-row/AMX K1b grouped planar IDOT), NOT a test gate. # Build on demand: make tests/bench_i4p_gidot ARCH=native -tests/bench_i4p_gidot$(EXE): tests/bench_i4p_gidot.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h idot.h sample.h kv_persist.h telemetry.h route_trace.h +tests/bench_i4p_gidot$(EXE): tests/bench_i4p_gidot.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h route_trace.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) # bench_gemv_stream: microbenchmark (decode-regime GEMV bandwidth vs the read ceiling; # frozen-baseline + deinterleaved-x candidate A/B), NOT a test gate. # Build on demand: make tests/bench_gemv_stream ARCH=native -tests/bench_gemv_stream$(EXE): tests/bench_gemv_stream.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h idot.h sample.h kv_persist.h telemetry.h route_trace.h +tests/bench_gemv_stream$(EXE): tests/bench_gemv_stream.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h route_trace.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) # bench_mla_simd: microbenchmark (scalar vs AVX2/NEON MLA-absorb reductions, #442), # NOT a test gate. Build on demand: make tests/bench_mla_simd -tests/bench_mla_simd$(EXE): tests/bench_mla_simd.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h idot.h sample.h kv_persist.h telemetry.h route_trace.h +tests/bench_mla_simd$(EXE): tests/bench_mla_simd.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h route_trace.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) -tests/test_uring$(EXE): tests/test_uring.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h idot.h sample.h kv_persist.h telemetry.h route_trace.h +tests/test_uring$(EXE): tests/test_uring.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h route_trace.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) tests/test_pipe_block$(EXE): tests/test_pipe_block.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) -tests/test_pilot_ring$(EXE): tests/test_pilot_ring.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h idot.h sample.h kv_persist.h telemetry.h +tests/test_pilot_ring$(EXE): tests/test_pilot_ring.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) -tests/test_moe_gs_guard$(EXE): tests/test_moe_gs_guard.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h idot.h sample.h kv_persist.h telemetry.h +tests/test_moe_gs_guard$(EXE): tests/test_moe_gs_guard.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) tests/test_omp_tune$(EXE): tests/test_omp_tune.c omp_tune.h compat.h @@ -2043,7 +2043,7 @@ tests/test_omp_tune$(EXE): tests/test_omp_tune.c omp_tune.h compat.h tests/test_compat_env$(EXE): tests/test_compat_env.c compat.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) -tests/test_kvb_notice$(EXE): tests/test_kvb_notice.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h idot.h sample.h kv_persist.h telemetry.h +tests/test_kvb_notice$(EXE): tests/test_kvb_notice.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) # Standalone: proves the group-scaled int8 GEMV keeps every row's float @@ -2216,7 +2216,7 @@ bench: iobench$(EXE) tests/test_kv_prefix$(EXE): tests/test_kv_prefix.c kv_prefix.h $(CC) $(CFLAGS) tests/test_kv_prefix.c -o tests/test_kv_prefix$(EXE) $(LDFLAGS) -tests/test_kimi_request_state$(EXE): tests/test_kimi_request_state.c kimi_k3.c kv_prefix.h serve_codec.h st.h json.h tok.h tok_unicode.h tok_unicode_o200k.h compat.h quant.h idot.h omp_tune.h route_trace.h +tests/test_kimi_request_state$(EXE): tests/test_kimi_request_state.c kimi_k3.c kv_prefix.h serve_codec.h st.h json.h tok.h tok_unicode.h tok_unicode_o200k.h compat.h quant.h fp8_format.h idot.h omp_tune.h route_trace.h $(CC) $(NOCUDA_CFLAGS) tests/test_kimi_request_state.c $(VK_OBJ) -o tests/test_kimi_request_state$(EXE) $(NOCUDA_LDFLAGS) # Kimi CUDA dispatch/fallback with a fake backend; no CUDA toolkit required. diff --git a/c/Makefile.deepseek-v4 b/c/Makefile.deepseek-v4 index a7858363f..5426cecda 100644 --- a/c/Makefile.deepseek-v4 +++ b/c/Makefile.deepseek-v4 @@ -246,20 +246,20 @@ $(REGISTRY_OBJ): expert_store_registry.c expert_store_registry.h expert_store.h $(CC) $(CFLAGS) -c expert_store_registry.c -o $@ $(TARGET_OBJS) $(TEST_UNIT_OBJS): %.o: deepseek_v4.c deepseek_v4.h \ - deepseek_v4_internal.h deepseek_v4_dspark.inc st.h json.h compat.h tensor.h quant.h \ + deepseek_v4_internal.h deepseek_v4_dspark.inc st.h json.h compat.h tensor.h quant.h fp8_format.h \ route_trace.h \ native_quant.h native_quant_batch.h native_quant_dual.h \ native_quant_fp4_rows16.h expert_store_registry.h $(CC) $(CFLAGS) -D$* -c deepseek_v4.c -o $@ $(V4_HOT_TEST_OBJ): deepseek_v4.c deepseek_v4.h deepseek_v4_internal.h \ - st.h json.h compat.h tensor.h quant.h route_trace.h \ + st.h json.h compat.h tensor.h quant.h fp8_format.h route_trace.h \ native_quant.h native_quant_fp4_rows16.h expert_store_registry.h $(CC) $(CFLAGS) -DCOLI_V4_TEST_HOOKS \ -DCOLI_V4_UNIT_EXPERT_STORE_HOT_ROWS16 -c deepseek_v4.c -o $@ $(V4_BATCH_TEST_OBJ): deepseek_v4.c deepseek_v4.h deepseek_v4_internal.h \ - tensor.h quant.h native_quant.h native_quant_batch.h + tensor.h quant.h fp8_format.h native_quant.h native_quant_batch.h $(CC) $(CFLAGS) -DCOLI_V4_TEST_HOOKS \ -DCOLI_V4_UNIT_NATIVE_QUANT_BATCH -c deepseek_v4.c -o $@ diff --git a/c/backend_metal.mm b/c/backend_metal.mm index e3497320c..3b767e2b3 100644 --- a/c/backend_metal.mm +++ b/c/backend_metal.mm @@ -767,7 +767,7 @@ static size_t fmt_bytes(int fmt, int I, int O) { // Grouped-int4 (fmt=4) scale-array size: one f32 per gsz-element group, per row -> O*ceil(I/gsz). // fp8 (fmt=8) scale-array size: one f32 per 128x128 BLOCK -> ceil(O/128)*ceil(I/128) (2D, // not per-row -- quant.h isn't included here, so the ceil-div is inlined rather than sharing -// colibri.c's qt_scale_bytes/quant.h's fp8_nblk). The block is a fixed 128x128, so gs is +// colibri.c's qt_scale_bytes/fp8_format.h's fp8_nblk). The block is a fixed 128x128, so gs is // ignored for fmt==8. f32 is this build's implemented scale // ENCODING for fmt=8 (see quant.h/colibri.c) -- this file has no reason to know that a // UE8M0 encoding exists at all: qt_resolve_fmt refuses it on the CPU read path before any diff --git a/c/colibri.c b/c/colibri.c index a96ccba4e..64bf36801 100644 --- a/c/colibri.c +++ b/c/colibri.c @@ -286,9 +286,10 @@ static int64_t qt_bytes(const QT *t){ /* byte residenti del tensore */ return (int64_t)t->O*(((int64_t)t->I+255)/256)*98 + 4; if(t->fmt==8){ /* fp8-e4m3 passthrough: O*I raw e4m3 bytes (n, byte-identical layout * to fmt=1's weight bytes) + one f32 scale per 128x128 block - * (FP8_BLOCK in quant.h, included below qt_bytes -- keep the - * arithmetic literal here, same discipline as fmt=5's comment - * above). Missing this branch would fall through to the fmt=2 + * (FP8_BLOCK in fp8_format.h via quant.h, included below + * qt_bytes -- keep the arithmetic literal here, same + * discipline as fmt=5's comment above). Missing this branch would + * fall through to the fmt=2 * default below (packed-nibble formula, ~half the real weight * bytes) and undercount a resident fp8 tensor's byte footprint -- * feeds AUTOPIN/RAM-budget math, so this branch is load-bearing diff --git a/c/fp8_format.h b/c/fp8_format.h new file mode 100644 index 000000000..4c465c2b0 --- /dev/null +++ b/c/fp8_format.h @@ -0,0 +1,30 @@ +/* fmt=8 (fp8-e4m3-b128) block geometry -- the single definition site for the + * 128x128 scale-block edge. Every consumer reaches these symbols through + * quant.h, which includes this header where the definitions used to live: + * matmul_fp8 and the e4m3 dequant plumbing (quant.h), the qt_addrow / + * qt_matvec_rows / qt_from_disk paths (colibri.c), and the native-FP8 + * reduction (qwen38.c / qwen38_core.h). Kept deliberately tiny (no LUTs, no + * functions with OpenMP pragmas, no intrinsics) so a translation unit that + * cannot include quant.h itself can still name the block edge instead of + * restating 128 / >>7 as literals. + * + * The in-repo FP8 tests are NOT independent of this constant: the reference + * decoders in test_fp8_passthrough.c, test_fp8_load.c, test_fp8_e2e_loader.c + * and test_qwen38_native_weights.c all index with FP8_BLOCK/fp8_nblk via + * quant.h, so they move in lockstep with an edit here (only + * test_backend_metal.mm's ref_fp8_nblk keeps its own literal, on purpose). + * An edit to FP8_BLOCK is a format change, not a tunable -- those tests + * cannot catch a wrong edit on their own. */ +#ifndef COLI_FP8_FORMAT_H +#define COLI_FP8_FORMAT_H + +#include + +#define FP8_BLOCK 128 + +/* Blocks covering n elements: ceil(n/FP8_BLOCK). Plain C on purpose -- + * nothing in this header may need decoration or a special front end, so it + * stays includable from any translation unit. */ +static inline int64_t fp8_nblk(int n){ return ((int64_t)n + FP8_BLOCK - 1) / FP8_BLOCK; } + +#endif /* COLI_FP8_FORMAT_H */ diff --git a/c/quant.h b/c/quant.h index 91cf75493..2c6486781 100644 --- a/c/quant.h +++ b/c/quant.h @@ -518,8 +518,10 @@ static inline __m256 bf16_decode8(const uint16_t *p) { } #endif -#define FP8_BLOCK 128 -static inline int64_t fp8_nblk(int n){ return ((int64_t)n + FP8_BLOCK - 1) / FP8_BLOCK; } +/* FP8_BLOCK / fp8_nblk moved to fp8_format.h so the CUDA backend shares the + * same named constant instead of restating 128 as literals (see that header's + * comment for the drift hazard this closes). */ +#include "fp8_format.h" /* y[S,O] = x[S,I] @ W^T, W raw e4m3 bytes (byte-identical layout to fmt=1) + * per-128x128-BLOCK f32 scale [ceil(O/128),ceil(I/128)]. Scalar reference path diff --git a/c/tests/test_makefile_deps.py b/c/tests/test_makefile_deps.py index be1c48e56..cfb2903a7 100644 --- a/c/tests/test_makefile_deps.py +++ b/c/tests/test_makefile_deps.py @@ -84,6 +84,58 @@ def test_the_engine_list_is_derived_and_not_empty(self): f"{family.id}: registered family whose build target is not " f"a checkable `{family.build_target}$(EXE):` rule") + def test_every_engine_rule_lists_fp8_format_h_if_it_reaches_it(self): + """A prerequisite reached THROUGH another header is still a prerequisite. + + The check below walks only an engine's own `#include "..."` lines, so a + header pulled in transitively is invisible to it. fp8_format.h is + exactly that shape: quant.h includes it, no engine .c does, so the + direct check stays green whether or not a rule lists it -- while + `touch fp8_format.h; make colibri` would report success and leave a + stale binary carrying the previous FP8_BLOCK. That is the drift the + header's own comment says it exists to prevent, so it is pinned here. + + Bite: delete `fp8_format.h` from `colibri$(EXE):` and this fails naming + colibri; the direct check above still passes. + + Scoped to fp8_format.h ON PURPOSE. Running the same closure over every + header reports seven rules that predate this change and are unrelated + to it -- colibri and deepseek_v41 do not list tok_unicode_o200k.h, + five engines do not list decode_batch.h or edge_adapter_internal.h, + qwen36 does not list sse41_kernels.h. Those are real instances of the + same hazard and worth a separate fix; widening this test to cover them + here would make it fail on arrival for reasons this change did not + cause. + """ + problems = [] + for target, (source, prereqs) in sorted(_engine_rules().items()): + seen, stack = set(), [source] + while stack: + cur = stack.pop() + try: + text = cur.read_text(encoding="utf-8") + except (OSError, UnicodeDecodeError): + continue + for h in INCLUDE_RE.findall(text): + if h in seen: + continue + path = C_DIR / h + # Only headers that exist in c/ are ours to track; system + # headers and generated files are not prerequisites. + if path.exists(): + seen.add(h) + stack.append(path) + if "fp8_format.h" in seen and "fp8_format.h" not in prereqs: + problems.append(f"{target} ({source.name}) reaches fp8_format.h " + f"through its include graph but does not list it") + self.assertEqual( + problems, [], + "fp8_format.h is reachable from these engines but absent from their " + "Makefile prerequisites. Editing it alone will NOT relink them -- " + "make reports success and leaves a stale artifact:\n " + + "\n ".join(problems), + ) + def test_every_engine_rule_lists_the_headers_its_source_includes(self): problems = [] for target, (source, prereqs) in sorted(_engine_rules().items()): diff --git a/docs/FORMATS.md b/docs/FORMATS.md index c7724d8aa..b52e90c8d 100644 --- a/docs/FORMATS.md +++ b/docs/FORMATS.md @@ -104,29 +104,29 @@ allowlist, :3379), consulted by both gate sites — `attention_rows` (:3448) and Sources for all rows (`c/quant.h`/`c/colibri.c` line numbers at this PR pair's current restack, base dev `292ed4c`): -- **fmt=0/1/2/3** — allocation policy: `qt_alloc`, `c/colibri.c:1105` +- **fmt=0/1/2/3** — allocation policy: `qt_alloc`, `c/colibri.c` (`bits>=16→fmt=0`, `bits>=5→fmt=1`, `bits>=4→fmt=2`, else `fmt=3`). - Kernels: `matmul_q` (`quant.h:105`, fmt=1), `matmul_i4` (`quant.h:125`, - fmt=2), `matmul_i2` (`quant.h:251`, fmt=3); pack/quantize helpers - `quantize_rows` (`quant.h:928`, fmt=1) and `pack_int2` (`quant.h:980`, - fmt=3). Byte-count formulas: `qt_bytes`, `c/colibri.c:183`. -- **fmt=4** (`int4-grouped`) — kernel `matmul_i4_grouped`, `quant.h:168`; + Kernels: `matmul_q` (`quant.h`, fmt=1), `matmul_i4` (`quant.h`, + fmt=2), `matmul_i2` (`quant.h`, fmt=3); pack/quantize helpers + `quantize_rows` (`quant.h`, fmt=1) and `pack_int2` (`quant.h`, + fmt=3). Byte-count formulas: `qt_bytes`, `c/colibri.c`. +- **fmt=4** (`int4-grouped`) — kernel `matmul_i4_grouped`, `quant.h`; group size `gs` is per-tensor, not fixed at 64 (contrast fmt=5). Byte-count: - `qt_bytes`'s `fmt==4` branch (inside `c/colibri.c:183`); scale-count split: - `qt_scale_bytes`, `c/colibri.c:263`. + `qt_bytes`'s `fmt==4` branch (inside `c/colibri.c`); scale-count split: + `qt_scale_bytes`, `c/colibri.c`. - **fmt=5** (`int3-g64`) — group size is fixed (`I3_GROUP=64`, - `quant.h:293`; `I3_GBYTES=24`, `quant.h:294`); helpers `i3_groups` - (`quant.h:295`), `i3_rowbytes` (`quant.h:296`); kernel `matmul_i3` - (`quant.h:354`); pack helper `pack_int3_g64` (`quant.h:956`). Allocation: - `qt_alloc`'s `bits==3` branch (inside `c/colibri.c:1105`). + `quant.h`; `I3_GBYTES=24`, `quant.h`); helpers `i3_groups` + (`quant.h`), `i3_rowbytes` (`quant.h`); kernel `matmul_i3` + (`quant.h`); pack helper `pack_int3_g64` (`quant.h`). Allocation: + `qt_alloc`'s `bits==3` branch (inside `c/colibri.c`). - **fmt=6** (`e8-iq3-lattice`) — upstream's merged code: format section header - precedes `quant.h:1008`; constants `E8_QK=256` (`quant.h:1008`), - `E8_SUB=32` (`quant.h:1009`), `E8_BBYTES=98` (`quant.h:1010`); row-byte - helpers `e8_blocks`/`e8_rowbytes` (`quant.h:1011-1012`); rotation contract - documented at `quant.h:1305` ("fmt=6 stores W@Q, so activations must be + precedes `quant.h`; constants `E8_QK=256` (`quant.h`), + `E8_SUB=32` (`quant.h`), `E8_BBYTES=98` (`quant.h`); row-byte + helpers `e8_blocks`/`e8_rowbytes` (`quant.h`); rotation contract + documented at `quant.h` ("fmt=6 stores W@Q, so activations must be transformed before"). Loader discriminator, upstream form (dev, ns==4 tag check at the top of `qt_resolve_fmt`): this branch's SECOND DESIGN - LANDMINE comment (`qt_resolve_fmt`, `c/colibri.c:1356`) hardens that check + LANDMINE comment (`qt_resolve_fmt`, `c/colibri.c`) hardens that check against the degenerate collisions below without changing any genuine-fmt=6 outcome. - **fmt=7** (`mxfp4`, upstream's merged code) — Vulkan-only decode: shader @@ -139,27 +139,27 @@ pair's current restack, base dev `292ed4c`): CPU (`quant.h`) or Metal kernel exists for it, and `qt_resolve_fmt` has no byte-arithmetic branch that returns 7. - **fmt=8** (`fp8-e4m3-b128`, this branch) — decode table `E4M3_LUT` - (`quant.h:446`) / `e4m3_decode` (`quant.h:480`), block size - `FP8_BLOCK=128` (`quant.h:482`), kernel `matmul_fp8` (`quant.h:491`). + (`quant.h`) / `e4m3_decode` (`quant.h`), block size + `FP8_BLOCK=128` (`fp8_format.h`), kernel `matmul_fp8` (`quant.h`). Disambiguation from fmt=1 ("THE DESIGN LANDMINE" — the two formats' weight bytes are byte-identical and can only be told apart by scale-array geometry, which is ambiguous for some small shapes) and the fmt=6 collision ("SECOND DESIGN LANDMINE") both live in `qt_resolve_fmt` - (`c/colibri.c:1356`), which now also consults an optional `stamped_name` + (`c/colibri.c`), which now also consults an optional `stamped_name` parameter (this PR): for the fmt=6 collision, a stamp resolves what an absent stamp still refuses; for the fmt=1-vs-fmt=8 collision, an absent stamp already resolves to `int8-row` since the #528 INVERSION, and a stamp's role there is instead letting a genuinely-stamped `fmt=8` tensor override that default — see "The metadata stamp" below for the exact rule in both cases. FMT_NAMES table (`name string` to `fmt int`): - `c/colibri.c:1316`. + `c/colibri.c`. - **no ordinal** (`int4-rans256-g0`, merged tools-only tier — line numbers at dev `7fb1159`, post-#671 merge `a3a5a75`, not at this PR pair's restack base) — codec + record reader/writer: `c/rans.h` - (`RANS_NSTREAMS 256`, `c/rans.h:93`; the record layout in the file-header - comment, `c/rans.h:17-34`; that same header names its engine consumer - "a future engine decode stage", `c/rans.h:4` — the format's own statement - that none exists yet). Identity constants: `c/tools/rans_format.py:40-46` + (`RANS_NSTREAMS 256`, `c/rans.h`; the record layout in the file-header + comment, `c/rans.h`; that same header names its engine consumer + "a future engine decode stage", `c/rans.h` — the format's own statement + that none exists yet). Identity constants: `c/tools/rans_format.py` (`FORMAT_NAME = "int4-rans256-g0"`; `METADATA_KEY = "colibri.fmt"` — the same key/shape as the fmt=8 stamp convention below, but MANDATORY here rather than a cross-check, because there is no byte arithmetic to fall @@ -169,15 +169,15 @@ pair's current restack, base dev `292ed4c`): named refusal classes in its module docstring). Full specification: `docs/int4-rans256-g0.md`. Engine-interaction status, stated precisely (why the ordinal column is empty, and what still runs): `qt_resolve_fmt` - (`c/colibri.c:1374`) has no branch that returns this format, and + (`c/colibri.c`) has no branch that returns this format, and `c/colibri.c`/`c/quant.h`/`c/st.h` contain no reference to it — no decode path, hence no ordinal. But a repacked shard is not invisible to - the engine: `st_fmt_stamp_ingest` (`c/st.h:317`, called from + the engine: `st_fmt_stamp_ingest` (`c/st.h`, called from `st_init_multi`'s discovery loop) parses its mandatory `colibri.fmt` stamp map at container-discovery time, and the three routed-expert load sites — exactly this format's target population — resolve formats by - byte arithmetic alone with `stamped_name=NULL` (`c/colibri.c:2217`, - `c/colibri.c:2386`, `c/colibri.c:2580`, each marked + byte arithmetic alone with `stamped_name=NULL` (`c/colibri.c`, + `c/colibri.c`, `c/colibri.c`, each marked `/* routed expert: never stamped */`), so the stamp is never consulted where it would matter most. A future consumer must wire stamp-gated dispatch AHEAD of that inference — `docs/int4-rans256-g0.md`'s From 0ad3eef770669786fa568a60e9d1e90efb409450 Mon Sep 17 00:00:00 2001 From: monotophic Date: Thu, 3 Sep 2026 13:50:58 -0400 Subject: [PATCH 2/6] feat(colibri): add fmt=8 arms to the CPU absorb path qt_addrow and qt_matvec_rows had no fmt=8 case, so an fp8-e4m3 kv_b_proj fell through to the per-row-scale tail, read t->s[row] past the end of a block-scale array and dereferenced a NULL q4. Both now decode fp8 directly, taking the block scale from [ceil(O/128), ceil(I/128)] exactly as matmul_fp8 does. qt_matvec_rows' arm mirrors that kernel statement for statement, including the f32-partial-per-block, double-across-blocks accumulation. qt_addrow is an axpy, so it mirrors the block-scale indexing only and folds the coefficient into the scale, the convention its fmt=4 arm has always used. A NaN block scale is refused by name, with the block that carried it: it poisons the whole block and every accumulator downstream of it, and these functions cannot repair it. A ZERO block scale is not refused -- it is valid data. Block scales are amax/448, so an all-zero block legitimately carries one, and decoding it as zeros is the correct answer; unlike on the GPU there is no table here that could be unwritten, because the CPU decoder reads e4m3 from a compile-time constant. The check costs one branch per block, not per element, and both properties are pinned: a NaN scale must refuse through either function, a zero scale must decode to zeros through both. The format guard below the new arms keeps its condition; only its message changes, to say that fmt=8 now returns above it. layer_cuda_shard_kvb gains a format allowlist. fmt 5/6/7 previously computed an int4 row stride against a group-scaled layout, took a non-NULL q4 and uploaded garbage silently; they are now refused by name before any pointer or stride is used, with a notice bounded to one line per format per process. --- c/colibri.c | 129 ++++++++++++--- c/fp8_format.h | 8 +- c/tests/test_qt_addrow.c | 259 ++++++++++++++++++++++++------ c/tools/repack_fp8_passthrough.py | 29 ++-- 4 files changed, 338 insertions(+), 87 deletions(-) diff --git a/c/colibri.c b/c/colibri.c index 64bf36801..8ffd06bde 100644 --- a/c/colibri.c +++ b/c/colibri.c @@ -2266,6 +2266,43 @@ static void qt_cuda_colocate(QT *dst,const QT *src){ } static void layer_cuda_shard_kvb(Layer *l,int H,int Q,int V){ if(!g_cuda_enabled||!g_cuda_dense||g_cuda_ndev<2||l->kv_b.fmt==0)return; + /* SHARD FORMAT ALLOWLIST (explicit refusal; this was an ACCIDENTAL fail-safe): the + * rb/weights/scale arithmetic below is written for exactly fmt=1 (int8, per-row + * scale), fmt=2 (int4 per-row), fmt=3 (int2 per-row) and fmt=4 (int4 grouped). + * Any other fmt reaching it computes a wrong row-byte stride, takes l->kv_b.q4 as + * the weight pointer (NULL for fmt=8, whose raw e4m3 bytes live in q8 -- see the + * QT struct comment), and slices l->kv_b.s with per-row/per-group geometry that + * fmt=8's per-128x128-BLOCK scales (and fmt=6's single 4-byte tag) simply do not + * have. fmt=8 only ever "worked" here by accident: q4==NULL made + * coli_cuda_tensor_upload_g's !weights check reject the upload before anything + * dereferenced it -- silent, unnamed, and one refactor away from a misread. + * Refuse BY NAME instead, BEFORE any pointer/stride use, and say what happens + * instead: the un-sharded kv_b stays whole on its layer home device, where fmt=8 + * kv_b decode runs the CPU absorb path (qt_addrow/qt_matvec_rows' fmt=8 + * branches below; the CUDA absorb kernels refuse fmt=8 via absorb_fmt_ok's + * fmt 0..4 allowlist) -- COLI_CUDA_ATTN_SHARD is a no-op for it. Same + * "refuse rather than misread" discipline as qt_addrow/ + * qt_matvec_rows' guards; notice only (no exit): sharding is an opt-in + * optimization and skipping it is the correct, working behavior. Bounded once + * per process per fmt, never per layer (metal_fmt_gate_notice, the precedent + * for bounded notices, is coarser still: one line per tensor KIND, naming only + * the first offending fmt). */ + if(l->kv_b.fmt!=1&&l->kv_b.fmt!=2&&l->kv_b.fmt!=3&&l->kv_b.fmt!=4){ + static int refused_fmt[32]; + if(!refused_fmt[l->kv_b.fmt&31]){ refused_fmt[l->kv_b.fmt&31]=1; + if(l->kv_b.fmt==8) + fprintf(stderr,"layer_cuda_shard_kvb: kv_b fmt=8 (fp8-e4m3, per-128x128-block " + "scales) has no head-shard layout here -- refusing the shard; fmt=8 kv_b " + "runs the absorb path on the layer home device instead, so " + "COLI_CUDA_ATTN_SHARD is a no-op for it (applies to every layer)\n"); + else + fprintf(stderr,"layer_cuda_shard_kvb: unsupported kv_b fmt=%d for the head-shard " + "upload (only fmt 1/2/3/4 match the per-row byte/scale strides computed " + "here) -- refusing the shard; kv_b stays whole on its layer home device " + "(applies to every layer)\n",l->kv_b.fmt); + } + return; + } int rb=l->kv_b.fmt==1?l->kv_b.I: (l->kv_b.fmt==2||l->kv_b.fmt==4)?(l->kv_b.I+1)/2:(l->kv_b.I+3)/4; const uint8_t *weights=l->kv_b.fmt==1?(const uint8_t*)l->kv_b.q8:l->kv_b.q4; @@ -3858,6 +3895,32 @@ static void expert_prefetch(Model *m, int layer, int eid){ /* ---- helper per l'ABSORPTION: accesso per-riga ai QT quantizzati ---- */ /* acc[0..I) += coef * W[row,:] (dequant al volo) */ +/* One fmt=8 block scale, checked before it multiplies anything. + * + * A NaN scale has no safe interpretation: it poisons the whole block and every + * accumulator downstream of it, and this function cannot repair it, so it is + * refused by name with the block that carried it -- the same "refuse rather + * than misread" discipline the format guards below apply. + * + * A ZERO scale is NOT refused: it is valid data. Block scales are amax/448, so + * a genuinely all-zero block (padding, an unused slice) legitimately produces + * zero, and decoding it as zeros is the correct answer. It cannot be confused + * with a decode against an unwritten table the way it can on the GPU, because + * there is no table here to be unwritten -- the CPU decoder reads its e4m3 + * values from a compile-time constant table, which is why the LUT-ready gate + * exists only on the CUDA side. Do not re-add a zero refusal: it would reject + * valid checkpoints. + * + * The check is per BLOCK, not per element: one branch per FP8_BLOCK columns. */ +static float fp8_block_scale(float sc, int64_t blkO, int64_t bi, const char *who){ + if(isnan(sc)){ + fprintf(stderr,"%s: fmt=8 scale block [%lld,%lld] is NaN -- refusing rather than " + "propagate it through the absorb accumulator\n",who,(long long)blkO,(long long)bi); + exit(1); + } + return sc; +} + static void qt_addrow(const QT *t, int row, float coef, float *acc){ int I=t->I; if(t->fmt==0){ const float *w=t->qf+(int64_t)row*I; for(int i=0;i>2]>>((k&3)*2))&3)|(((hi[k>>3]>>(k&7))&1)<<2); acc[base+k]+=cg*(float)((int)u-4); } } return; } + /* fmt=8 (fp8-e4m3-b128, absorb-path support added here): t->s holds ONE f32 scale + * per 128x128 BLOCK (ceil(O/128)*ceil(I/128) entries, block-row-major), not O, and + * t->q4 is NULL for this format -- raw e4m3 bytes live in t->q8 instead, same + * convention as fmt=1 (see the QT struct comment). Mirrors matmul_fp8's (quant.h) + * block-scale indexing exactly: blkO=row/FP8_BLOCK selects the scale row, then one + * scale per FP8_BLOCK-wide slice of I. This is the branch that used to be missing + * -- see the guard below's history note. */ + if(t->fmt==8){ const uint8_t *w=(const uint8_t*)t->q8+(int64_t)row*I; + int64_t nblkI=fp8_nblk(I), blkO=(int64_t)row/FP8_BLOCK; + const float *scl=t->s+blkO*nblkI; + for(int64_t bi=0; bi*FP8_BLOCKI) blen=I-base; + float sc=coef*fp8_block_scale(scl[bi],blkO,bi,"qt_addrow"); + for(int i=base;is[row]) followed by fmt=1 (int8, explicit * branch), fmt=2 (int4 packed, explicit branch), or the tail's own IMPLICIT fmt=3 * (int2 packed, the final unconditional block) -- there was no guard stopping any @@ -3891,20 +3969,16 @@ static void qt_addrow(const QT *t, int row, float coef, float *acc){ * a heap OVERREAD, and the untouched fall-through then misreads t->q4's real E8 * lattice bytes as int2-packed data (same bug SHAPE as #298's CUDA absorb-kernel * fix, and the same one this file's own fmt=4/5 branches above were added to - * dodge -- fmt=6 was simply missed). fmt=8 (fp8-e4m3-b128): t->s holds - * ceil(O/128)*ceil(I/128) per-block floats, not O -- t->s[row] overreads for - * row>=nblk (e.g. a [130,130] tensor has nblk=4, so every row past 3 already reads - * out of bounds), AND t->q4 is NULL for fmt=8 (raw bytes live in t->q8 instead, - * same convention as fmt=1 -- see the QT struct comment), so the fall-through's - * `t->q4+(int64_t)row*((I+3)/4)` dereferences NULL-plus-offset: SIGSEGV, - * reproduced (see the report's proof-of-bite transcript). Refuse loudly instead -- - * this function has no byte-count context of its own to validate against (it only - * ever sees an already-resolved QT), so "unsupported fmt" is the only check - * available, same "refuse rather than misread" discipline qt_resolve_fmt applies - * at load time. */ + * dodge -- fmt=6 was simply missed). fmt=8 previously landed here too (t->s[row] + * overreads past row>=nblk, t->q4 is NULL -> SIGSEGV via the int2 fall-through); + * it now returns above via its own branch and never reaches this guard. Refuse + * loudly instead for anything else -- this function has no byte-count context of + * its own to validate against (it only ever sees an already-resolved QT), so + * "unsupported fmt" is the only check available, same "refuse rather than + * misread" discipline qt_resolve_fmt applies at load time. */ if(t->fmt!=1 && t->fmt!=2 && t->fmt!=3){ fprintf(stderr,"qt_addrow: unsupported fmt=%d for the per-row-scale absorb path " - "(only fmt 1/2/3 reach this point; fmt 0/4/5 are handled above and return " + "(only fmt 1/2/3 reach this point; fmt 0/4/5/8 are handled above and return " "before it) -- refusing rather than misread t->s[row]/t->q4\n", t->fmt); exit(1); } @@ -3954,18 +4028,35 @@ static void qt_matvec_rows(const QT *t, int r0, int n, const float *x, float *y) for(int k=0;k>2]>>((k&3)*2))&3)|(((hi[k>>3]>>(k&7))&1)<<2); acc+=(float)((int)u-4)*x[base+k]; } a+=(double)(acc*sr[g]); } } + /* fmt=8 (fp8-e4m3-b128, absorb-path support added here): per-128x128-BLOCK f32 + * scale, block-row-major (ceil(O/128)*ceil(I/128) entries), raw bytes in t->q8 + * (t->q4 is NULL for this format). Same block-scale indexing as matmul_fp8 + * (quant.h) and qt_addrow's fmt=8 branch above: blkO=row/FP8_BLOCK picks the + * scale row, one scale per FP8_BLOCK-wide slice of I, double-accumulated + * across blocks like matmul_fp8 to avoid unfairly penalizing cross-block + * cancellation (widen-then-multiply, a+=(double)acc*sc -- the same rounding + * as matmul_fp8 and this function's grouped fmt=4 arm; fmt=5's arm rounds + * differently, multiplying in float before widening). */ + else if(t->fmt==8){ const uint8_t *w=(const uint8_t*)t->q8+(int64_t)row*I; + int64_t nblkI=fp8_nblk(I), blkO=(int64_t)row/FP8_BLOCK; + const float *scl=t->s+blkO*nblkI; + for(int64_t bi=0; bi*FP8_BLOCKI) blen=I-base; + float sc=fp8_block_scale(scl[bi],blkO,bi,"qt_matvec_rows"); float acc=0; + for(int i=base;is is a fixed 4-byte tag (t->s[row] overreads for row>0), - * fmt=8's t->s holds per-128x128-block floats (t->s[row] overreads for - * row>=nblk) and t->q4 is NULL for fmt=8 -- both would have silently misread or - * crashed here exactly like qt_addrow did before its own fix; refuse instead. */ + * defect): fmt=6's t->s is a fixed 4-byte tag (t->s[row] overreads for row>0) -- + * it would silently misread or crash here exactly like qt_addrow did before its + * fix; refuse instead. fmt=8 previously fell into this same trap and now has its + * own branch above instead. */ else if(t->fmt==3){ const uint8_t *w=t->q4+(int64_t)row*((I+3)/4); float s=t->s[row]; float acc=0; for(int i=0;i>2]; acc+=((int)((b>>((i&3)*2))&3)-2)*x[i]; } a=acc*s; } else { fprintf(stderr,"qt_matvec_rows: unsupported fmt=%d for the per-row-scale absorb " - "path (only fmt 0/1/2/3/4/5 are handled) -- refusing rather than misread " + "path (only fmt 0/1/2/3/4/5/8 are handled) -- refusing rather than misread " "t->s[row]/t->q4\n", t->fmt); exit(1); } diff --git a/c/fp8_format.h b/c/fp8_format.h index 4c465c2b0..04610efef 100644 --- a/c/fp8_format.h +++ b/c/fp8_format.h @@ -9,10 +9,10 @@ * restating 128 / >>7 as literals. * * The in-repo FP8 tests are NOT independent of this constant: the reference - * decoders in test_fp8_passthrough.c, test_fp8_load.c, test_fp8_e2e_loader.c - * and test_qwen38_native_weights.c all index with FP8_BLOCK/fp8_nblk via - * quant.h, so they move in lockstep with an edit here (only - * test_backend_metal.mm's ref_fp8_nblk keeps its own literal, on purpose). + * decoders in test_fp8_passthrough.c, test_fp8_load.c, test_fp8_e2e_loader.c, + * test_qwen38_native_weights.c and test_qt_addrow.c all index with + * FP8_BLOCK/fp8_nblk via quant.h, so they move in lockstep with an edit here + * (only test_backend_metal.mm's ref_fp8_nblk keeps its own literal, on purpose). * An edit to FP8_BLOCK is a format change, not a tunable -- those tests * cannot catch a wrong edit on their own. */ #ifndef COLI_FP8_FORMAT_H diff --git a/c/tests/test_qt_addrow.c b/c/tests/test_qt_addrow.c index 46c3673ca..254c9e772 100644 --- a/c/tests/test_qt_addrow.c +++ b/c/tests/test_qt_addrow.c @@ -4,31 +4,35 @@ * fmt 0/4/5 explicitly, then fall through assuming a PER-ROW scale (t->s[row]) followed by * fmt=1/2/3 (qt_addrow) or fmt=0/1/2/3/4/5 (qt_matvec_rows, via an if/else-if chain ending * in a bare `else`) -- nothing stopped fmt=6 (E8/IQ3, t->s is a FIXED 4-byte tag, not O - * floats) or fmt=8 (fp8-e4m3-b128, t->s holds per-128x128-block floats, not O; t->q4 is - * NULL) from reaching that fall-through. For fmt=8 specifically this SIGSEGVs: t->s[row] - * overreads (silently, usually not fatal on its own), then the untouched tail computes - * `t->q4+(int64_t)row*((I+3)/4)` on a NULL t->q4 and dereferences it. For fmt=6 it silently - * misreads the real E8/IQ3 lattice bytes as int2-packed data (same bug SHAPE as #298's CUDA - * absorb-kernel fix, which is why this file's own fmt=4/5 branches exist -- fmt=6 was simply - * missed). Both functions now refuse loudly (exit(1), naming the function and the fmt) for - * any fmt they don't explicitly handle, matching qt_resolve_fmt's own "refuse rather than - * misread" discipline. + * floats) from reaching that fall-through, where it silently misreads the real E8/IQ3 + * lattice bytes as int2-packed data (same bug SHAPE as #298's CUDA absorb-kernel fix, which + * is why this file's own fmt=4/5 branches exist -- fmt=6 was simply missed). fmt=6 has no + * decoder in either function and still refuses loudly (exit(1), naming the function and the + * fmt), matching qt_resolve_fmt's own "refuse rather than misread" discipline. * - * This file: (1) proves the refusal fires for fmt=6 and fmt=8 through BOTH functions - * (fork+pipe+waitpid, this suite's established house pattern for exit(1)-terminated paths -- - * see tests/test_fp8_load.c's expect_refuse/expect_stamp_refuse); (2) proves every format - * BOTH functions still legitimately handle (0/1/2/3/4/5) produces byte-identical results - * against an independently-written reference dequantizer (qt_dequant_row_ref below -- NOT - * copy-pasted from qt_addrow/qt_matvec_rows, restructured as a single per-element loop per - * format, so a real regression in either the guard's placement or the untouched per-fmt math - * would show up here, not just tautologically re-run the same code). Reachability note (not - * a scope excuse, just context): both functions serve ONLY the kv_b absorb path - * (attention_rows/decode call sites), and tools/repack_fp8_passthrough.py deliberately - * excludes kv_b_proj from fmt=8 repacking -- so this fires only via a hand-slotted or - * ambiguous-collision container, not this repo's own tooling's own output. Crash-instead-of- - * refuse is still a real defect (spec I6: loud failure, every refusal names its condition), - * and the fmt=8 QT surface these functions can now be handed is one this same PR pair - * created. */ + * fmt=8 (fp8-e4m3-b128) USED to hit this same fall-through and SIGSEGV (t->s[row] overread, + * then a NULL t->q4 dereference -- see git history for the pre-fix account) but now has its + * own explicit branch in both functions (the absorb-decode commit this file accompanies): + * t->s holds per-128x128-block scales, t->q8 holds raw e4m3 bytes (t->q4 stays NULL), decoded + * through the same e4m3 LUT / block-scale geometry (quant.h's e4m3_decode/fp8_nblk/FP8_BLOCK) + * that matmul_fp8 already exercises on the dense/expert path. This file now (1) proves fmt=8 + * produces correct, tolerance-bounded output through BOTH functions -- an exact-dequant check + * against an independent per-element reference (qt_dequant_row_ref's new fmt==8 branch, + * disjoint code shape from qt_addrow/qt_matvec_rows' block-batched loops) plus a direct parity + * check against matmul_fp8 (the proven non-absorb fp8 reference) on a kv_b-shaped tensor; (2) + * proves the refusal still fires for fmt=6 through both functions (fork+pipe+waitpid, this + * suite's established house pattern for exit(1)-terminated paths -- see + * tests/test_fp8_load.c's expect_refuse/expect_stamp_refuse); (3) proves every other format + * both functions handle (0/1/2/3/4/5) produces byte-identical results against the same + * independent reference dequantizer, so a real regression in either the guard's placement or + * the untouched per-fmt math would show up here, not just tautologically re-run the same code. + * Reachability note (not a scope excuse, just context): both functions serve ONLY the kv_b + * absorb path (attention_rows/decode call sites), and the repo's own repack tool mints + * exactly the container these arms decode -- tools/repack_fp8_passthrough.py emits kv_b_proj + * (kind "kvb", in its RESIDENT_KINDS/STAMPABLE_KINDS sets) as byte-preserved fmt=8 with a + * stamped scale sidecar, and attention_rows' `int absorb = kvs || ...` makes absorb + * unbypassable on the batched serving path. These arms are the decode support that + * container needs. */ #define main coli_glm_main_unused #include "../colibri.c" #undef main @@ -48,13 +52,22 @@ static uint64_t rng = 0xA11CE5EEDF00Dull; static uint8_t rndbyte(void){ rng ^= rng << 13; rng ^= rng >> 7; rng ^= rng << 17; return (uint8_t)(rng & 0xFF); } static float rndsmallf(void){ rng ^= rng << 13; rng ^= rng >> 7; rng ^= rng << 17; return ((int64_t)(rng & 0xFFF) - 0x800) / (float)0x800; } /* [-1,1)-ish, small magnitude */ +/* fmt=8 fixtures must avoid the two NaN byte patterns (0x7F/0xFF, see quant.h's E4M3_LUT + * comment) -- a NaN weight is a real, policy-accepted outcome in production, but it would + * make this file's relative-error comparisons meaningless (NaN != NaN), same reasoning + * tests/test_fp8_load.c's rndbyte_nonan documents. */ +static uint8_t rndbyte_nonan(void){ + for(;;){ uint8_t b=rndbyte(); if(b!=0x7F && b!=0xFF) return b; } +} /* ---- independent reference dequantizer: W[row,:] as floats, one element per loop * iteration -- deliberately NOT the same code shape as qt_addrow/qt_matvec_rows (those * unroll pairs for fmt=2/3, split low/high planes for fmt=5) so this genuinely * cross-checks the production math, not just the production code running twice. ---- */ +static void qt_dequant_row_ref_fmt8(const QT *t, int row, float *out); static void qt_dequant_row_ref(const QT *t, int row, float *out){ int I=t->I; + if(t->fmt==8){ qt_dequant_row_ref_fmt8(t,row,out); return; } if(t->fmt==0){ const float *w=t->qf+(int64_t)row*I; for(int i=0;ifmt==4){ const uint8_t *w=t->q4+(int64_t)row*((I+1)/2); int gs=t->gs, ng=(I+gs-1)/gs; @@ -80,6 +93,22 @@ static void qt_dequant_row_ref(const QT *t, int row, float *out){ { const uint8_t *w=t->q4+(int64_t)row*((I+3)/4); for(int i=0;i>2]; int v=(b>>((i&3)*2))&3; out[i]=((int)v-2)*s; } } } +/* fmt=8 (fp8-e4m3-b128) independent reference: a flat per-ELEMENT loop, deliberately NOT + * the block-batched-accumulation shape qt_addrow/qt_matvec_rows and matmul_fp8 all share -- + * this recomputes each element's block index and looks up its scale independently, on every + * iteration, rather than walking block-by-block. Reuses quant.h's e4m3_decode/fp8_nblk/ + * FP8_BLOCK constants deliberately: those ARE the declared format geometry under test here + * (the same LUT/geometry every fmt=8 consumer in the tree shares), not implementation + * detail this reference should reinvent. */ +static void qt_dequant_row_ref_fmt8(const QT *t, int row, float *out){ + int I=t->I; + int64_t nblkI=fp8_nblk(I), blkO=(int64_t)row/FP8_BLOCK; + const float *scl=t->s+blkO*nblkI; + for(int i=0;iq8[(int64_t)row*I+i])*scl[bi]; + } +} /* ---- fixture builders: one per format, deterministic pseudo-random payload ---- */ static void fill_fmt0(QT *t, int O, int I){ @@ -126,6 +155,17 @@ static void fill_fmt5(QT *t, int O, int I){ for(int i=0;iq4[i]=rndbyte(); for(int i=0;is[i]=0.025f+0.0003f*(float)i; } +/* kv_b-shaped fmt=8 fixture: I is left caller-chosen so callers can pick both the + * real contraction dimension (kv_lora_rank=512, GLM-5.2's MLA latent width) and + * small tail-covering shapes. */ +static void fill_fmt8(QT *t, int O, int I){ + t->fmt=8; t->O=O; t->I=I; t->gs=0; + int64_t nblkO=fp8_nblk(O), nblkI=fp8_nblk(I), nblk=nblkO*nblkI; + t->q8=(int8_t*)malloc((size_t)O*I); + t->s=(float*)malloc((size_t)nblk*sizeof(float)); + for(int64_t i=0;i<(int64_t)O*I;i++) t->q8[i]=(int8_t)rndbyte_nonan(); + for(int64_t i=0;is[i]=0.01f+0.001f*(float)i; +} static void free_qt(QT *t){ free(t->qf); free(t->q8); free(t->q4); free(t->s); memset(t,0,sizeof *t); } /* ---- byte-identity: qt_addrow / qt_matvec_rows vs the independent reference, every @@ -185,22 +225,113 @@ static void test_byte_identity_all_formats(void){ memset(&t,0,sizeof t); fill_fmt3(&t,4,17); check_addrow_identity(&t,"fmt=3"); check_matvec_identity(&t,"fmt=3"); free_qt(&t); memset(&t,0,sizeof t); fill_fmt4(&t,4,40,16);check_addrow_identity(&t,"fmt=4"); check_matvec_identity(&t,"fmt=4"); free_qt(&t); memset(&t,0,sizeof t); fill_fmt5(&t,4,130); check_addrow_identity(&t,"fmt=5"); check_matvec_identity(&t,"fmt=5"); free_qt(&t); + /* fmt=8: small single-block-both-dims shape (mirrors the other formats' O=4-ish + * defaults) plus a kv_b-shaped multi-block shape (O=192 -> nblkO=2 with a partial + * tail block; I=512=kv_lora_rank -> nblkI=4, exact blocks) so both the partial-row- + * block and the multi-column-block paths through the new branches are exercised. */ + memset(&t,0,sizeof t); fill_fmt8(&t,6,17); check_addrow_identity(&t,"fmt=8 (single block)"); check_matvec_identity(&t,"fmt=8 (single block)"); free_qt(&t); + memset(&t,0,sizeof t); fill_fmt8(&t,192,512); check_addrow_identity(&t,"fmt=8 (kv_b-shaped, multi-block)"); check_matvec_identity(&t,"fmt=8 (kv_b-shaped, multi-block)"); free_qt(&t); + /* nblkI>=2 WITH a partial COLUMN tail: I=200 -> nblkI=2 with a 72-wide tail block, + * O=130 -> nblkO=2 with a 2-row tail. The kv_b-shaped case above has I=512 (exact + * column blocks), so the per-row scale STRIDE (nblkI) and the bi block-scale index + * were only ever exercised at shapes where flooring/misdeviating them is invisible. + * Mutation-checked: corrupting the block-scale index math (nblkI=I/FP8_BLOCK, the + * floor) passes every fmt=8 case above and fails exactly this one. */ + memset(&t,0,sizeof t); fill_fmt8(&t,130,200); check_addrow_identity(&t,"fmt=8 (partial column tail)"); check_matvec_identity(&t,"fmt=8 (partial column tail)"); free_qt(&t); +} + +/* ---- fmt=8 absorb-path vs the proven non-absorb matmul_fp8 reference (quant.h) ---- + * qt_matvec_rows(t,row,1,x,&y) and matmul_fp8's per-output-row computation are the SAME + * mathematical quantity (y[row] = sum_i x[i]*dequant(W[row,i])) computed with the exact + * same block-scale geometry and the exact same accumulation order (float-accumulate within + * a 128-wide block, double-accumulate the per-block partials) -- both were written to + * mirror that convention deliberately (see this PR's qt_matvec_rows fmt=8 branch comment + * in colibri.c). Compiled in the same translation unit with the same flags, so the two are + * expected to agree tightly; the tolerance below is a guard against incidental FP-contraction + * differences between the two call sites, not evidence of a real algorithmic mismatch. */ +static void test_fmt8_matmul_fp8_parity(void){ + enum { O=192, I=512 }; /* kv_b-shaped: I=kv_lora_rank=512, O spans a partial row-block */ + QT t; memset(&t,0,sizeof t); fill_fmt8(&t,O,I); + float *x=(float*)malloc((size_t)I*sizeof(float)); + for(int i=0;i1e-6f ? ae/fabsf(want) : ae; + if(rel > 1e-5f){ + printf("FAIL fmt=8 matmul_fp8 parity: row=%d got=%.9g want=%.9g rel=%.3g\n", + row,(double)y,(double)want,(double)rel); + fails++; + } + } + free(x); free(yref); free_qt(&t); +} + +/* ---- fmt=8 NaN propagation (policy pin -- see quant.h's "NaN POLICY" note): a NaN + * weight byte (0x7F/0xFF) decodes to a real IEEE NaN and is left to PROPAGATE, relying + * on the tested downstream sampler net (tests/test_logit_nan.c), never scrubbed at the + * weight level. The identity checks above deliberately exclude NaN bytes + * (rndbyte_nonan) -- and would pass NaN lanes silently anyway (NaN > eps compares + * false) -- so this pins the absorb-path behavior explicitly against the independent + * reference: qt_addrow poisons exactly the accumulator lanes whose reference dequant is + * NaN (per-element accumulate, both NaN byte codes), qt_matvec_rows poisons the whole + * dot product of any row containing one, and clean rows/lanes of the same tensor stay + * NaN-free and tolerance-identical to the reference. ---- */ +static void test_fmt8_nan_propagation(void){ + enum { O=130, I=200 }; /* same tail-covering shape as the identity case above */ + QT t; memset(&t,0,sizeof t); fill_fmt8(&t,O,I); + /* one NaN of EACH byte code, in different row-blocks and different column blocks; + * (129,199) lands in the tail row-block x tail column-block corner */ + const int nan_row[2]={3,129}, nan_col[2]={5,199}; + t.q8[(int64_t)nan_row[0]*I+nan_col[0]]=(int8_t)0x7F; + t.q8[(int64_t)nan_row[1]*I+nan_col[1]]=(int8_t)0xFF; + float *ref=(float*)malloc((size_t)I*sizeof(float)); + float *acc=(float*)malloc((size_t)I*sizeof(float)); + float *x=(float*)malloc((size_t)I*sizeof(float)); + for(int i=0;i1e-6f?ae/fabsf(want):ae; + if(rel>1e-5f){ printf("FAIL fmt=8 NaN: qt_addrow row=%d i=%d clean lane got=%.9g want=%.9g\n",row,i,(double)acc[i],(double)want); fails++; } + } + } + float y=0.f; qt_matvec_rows(&t,row,1,x,&y); + CHECK(isnan(y)); /* dot product over a NaN-bearing row is NaN, like the reference sum */ + } + /* a clean row of the SAME tensor stays NaN-free through both functions */ + { int row=4; float y=0.f; + qt_dequant_row_ref(&t,row,ref); + for(int i=0;is[0]=v; } + +static void call_addrow_fmt8_nan(void){ + QT t; memset(&t,0,sizeof t); fill_fmt8(&t,6,17); poison_scale(&t,(float)NAN); + float acc[17]; memset(acc,0,sizeof acc); + qt_addrow(&t,0,1.f,acc); +} +static void call_matvec_fmt8_nan(void){ + QT t; memset(&t,0,sizeof t); fill_fmt8(&t,6,17); poison_scale(&t,(float)NAN); + static float x[17]; for(int i=0;i<17;i++) x[i]=rndsmallf(); + float y=0.f; qt_matvec_rows(&t,0,1,x,&y); +} +/* A ZERO block scale is valid data, not corruption: block scales are amax/448, + * so an all-zero block legitimately carries one. Both arms must decode it to + * zeros in place -- not refuse it, and not leave the accumulator untouched. + * These two assertions exist to keep a zero refusal from being re-added. */ +static void test_fmt8_zero_scale_decodes_to_zeros(void){ + enum { O=6, I=17 }; + { QT t; memset(&t,0,sizeof t); fill_fmt8(&t,O,I); poison_scale(&t,0.f); + float acc[I]; for(int i=0;ikv_b) and the CUDA absorb kernels have no fmt==8 case as of this tool's own -HEAD -- that support is being built in parallel on `f8/absorb-fmt8` and is -NOT part of this diff. A container minted with this tool BEFORE that engine -branch lands will load clean (qt_from_disk resolves the stamp exactly like -o_proj's) but crash -- loudly, not silently -- on any decode that reaches the -batched (`kvs`-nonNULL) serving path, because `ABSORB=0` cannot bypass absorb -there. Sequencing that correctly (engine work before this container is used -for decode) is the gate's job, not this tool's: repacking kv_b_proj here is -the tool-side half of a two-sided integration, done now because engine and -tool work can (and per the gate's assembly manifest, should) proceed in -parallel. +never at the ambiguous [O,I] this tool's `_check_geometry` guards). The +engine-side half of this two-sided integration now exists on the CPU: +colibri.c's MLA-absorption path (qt_addrow/qt_matvec_rows, called only on +l->kv_b) carries explicit fmt==8 decode arms, so a container minted here +loads clean (qt_from_disk resolves the stamp exactly like o_proj's) and +decodes correctly through the batched (`kvs`-nonNULL) serving path, where +absorb cannot be bypassed. The CUDA absorb kernels still have no fmt==8 +case (absorb_fmt_ok delegates to coli_cuda_weight_at_supported's fmt 0..4 +allowlist), so fmt=8 kv_b decode stays on the CPU absorb path until that +lands. METADATA STAMP (reference implementation of the FORMATS-registry FR -- see docs/FORMATS.md): every output shard's safetensors `__metadata__` carries a From 9c7ae3871ce74a9235369eabbeb1d3b01753c5b7 Mon Sep 17 00:00:00 2001 From: monotophic Date: Thu, 3 Sep 2026 14:22:39 -0400 Subject: [PATCH 3/6] feat(cuda): add fmt=8 (fp8-e4m3) decode to the absorb kernels weight_at and absorb_scale gain fmt=8 branches so the attention absorb path decodes fp8 on the GPU, with the block-scale geometry of the dense fmt=8 matmul and of matmul_fp8. The surviving 128 literals are pinned to the shared constant by a file-scope assertion. The e4m3 table is per-device while the flag that gates fmt=8 uploads is process-wide, so the two must not drift. coli_cuda_shutdown clears the flag; coli_cuda_init never writes it. That is sufficient, and this states the mechanism rather than asserting the conclusion: init will not rebuild contexts underneath a live device set. A re-init naming the same set returns success and leaves the contexts -- and the table published to them -- untouched; one naming a different set is refused before anything is rebuilt. So the set cannot widen past what the last publish covered without passing through shutdown. The two decisions that gate rests on are factored into pure predicates in backend_cuda.h, which the backend calls at the real sites rather than keeping a second copy. tests/test_cuda_lut_gate.c pins them, and the resulting state machine, on a plain CPU build with no GPU and no CUDA toolchain: previously nothing about this gate ran anywhere except on a CUDA host, which is how three sites in the tree came to assert init semantics that no longer held. The lifecycle test in tests/test_fp8_cuda.cu now asserts all three edges, including that a same-set re-init keeps the flag -- requiring a republish there would force one that buys nothing. absorb_scale's own contract comment named fmt=4 as the only format without a per-row scale; it now names fmt=8's per-block layout as the other exception, since this change adds that branch inside the same function. --- c/Makefile | 26 +++-- c/backend_cuda.cu | 121 +++++++++++++++++------ c/backend_cuda.h | 62 ++++++++++-- c/colibri.c | 7 +- c/fp8_format.h | 37 ++++--- c/tests/test_backend_cuda.cu | 114 ++++++++++++++++++++++ c/tests/test_cuda_fmt_guard.c | 18 ++-- c/tests/test_cuda_fmt_trap_cuda.cu | 109 +++++++++++++++++++-- c/tests/test_cuda_lut_gate.c | 142 +++++++++++++++++++++++++++ c/tests/test_fp8_cuda.cu | 50 ++++++++++ c/tests/test_shard_kvb_refuse.c | 151 +++++++++++++++++++++++++++++ c/tools/repack_fp8_passthrough.py | 20 ++-- 12 files changed, 772 insertions(+), 85 deletions(-) create mode 100644 c/tests/test_cuda_lut_gate.c create mode 100644 c/tests/test_shard_kvb_refuse.c diff --git a/c/Makefile b/c/Makefile index 9fedfa799..1ee880c11 100644 --- a/c/Makefile +++ b/c/Makefile @@ -784,7 +784,7 @@ backend_loader.o: backend_loader.c backend_cuda.h compat.h .build-config # shell that has the MSVC environment set (e.g. after vcvars64.bat, or from a # "x64 Native Tools Command Prompt"). COLI_CUDA_BUILDING_DLL enables # __declspec(dllexport) so the 15 API symbols are exported. -cuda-dll: backend_cuda.cu backend_cuda.h +cuda-dll: backend_cuda.cu backend_cuda.h fp8_format.h @command -v "$(NVCC)" >/dev/null 2>&1 || { echo "nvcc not found: set CUDA_HOME or NVCC" >&2; exit 1; } @command -v cl >/dev/null 2>&1 || { echo "cl.exe (MSVC) not in PATH — run vcvars64.bat first" >&2; exit 1; } # The banner is localized ("for x64" / "per x64" / "pour x64"): match the arch token only (#1531). @@ -814,7 +814,7 @@ cuda-dll: backend_cuda.cu backend_cuda.h # survives. No path is hardcoded — see the HIP SDK contract above. # # make hip-dll HIP_DLL=1 HIP_SDK_ROOT= HIP_ARCH=gfx1151 -hip-dll: backend_cuda.cu backend_cuda.h backend_gpu_compat.h +hip-dll: backend_cuda.cu backend_cuda.h fp8_format.h backend_gpu_compat.h @test -x "$(HIPCC)" || command -v "$(HIPCC)" >/dev/null 2>&1 || { echo "hipcc not found at \"$(HIPCC)\": set HIP_SDK_ROOT=, HIP_BIN_DIR= or HIPCC=" >&2; exit 1; } @test -d "$(HIP_INCLUDE_DIR)" || { echo "HIP include dir not found: \"$(HIP_INCLUDE_DIR)\" — set HIP_INCLUDE_DIR=" >&2; exit 1; } @test -f "$(HIP_INCLUDE_DIR)/hip/hip_runtime.h" || { echo "hip/hip_runtime.h missing under \"$(HIP_INCLUDE_DIR)\" — set HIP_INCLUDE_DIR=" >&2; exit 1; } @@ -833,7 +833,7 @@ backend_cuda_ink.o: backend_cuda_ink.cu backend_cuda_ink.h .build-config @command -v "$(NVCC)" >/dev/null 2>&1 || { echo "nvcc not found: set CUDA_HOME or NVCC" >&2; exit 1; } "$(NVCC)" $(NVCCFLAGS) -c backend_cuda_ink.cu -o $@ -backend_cuda.o: backend_cuda.cu backend_cuda.h backend_gpu_compat.h .build-config +backend_cuda.o: backend_cuda.cu backend_cuda.h fp8_format.h backend_gpu_compat.h .build-config @command -v "$(GPUCC)" >/dev/null 2>&1 || { echo "$(GPUCC_NAME) not found" >&2; exit 1; } "$(GPUCC)" $(GPUFLAGS) -c backend_cuda.cu -o $@ @@ -956,7 +956,7 @@ rans: $(RANSLIB) $(RANSLIB): tools/rans_ctypes.c rans.h $(CC) $(CFLAGS) -fPIC -shared $< -o $@ $(LDFLAGS) -cuda-test: tests/test_mxfp4_expert_cuda.cu backend_cuda.cu backend_cuda.h backend_gpu_compat.h tests/test_backend_cuda.cu tests/test_ragged_attention.cu tests/test_absorb_determinism.cu tests/test_mxfp4_cuda.cu tests/mxfp4_ref.c tests/test_fp8_warp_cuda.cu tests/test_fp8_cuda.cu tests/test_weights_owned_cuda.cu tests/test_cuda_fmt_trap_cuda.cu tests/test_alloc_footprint_cuda.cu tests/test_cuda_init_failure.cu +cuda-test: tests/test_mxfp4_expert_cuda.cu backend_cuda.cu backend_cuda.h fp8_format.h backend_gpu_compat.h tests/test_backend_cuda.cu tests/test_ragged_attention.cu tests/test_absorb_determinism.cu tests/test_mxfp4_cuda.cu tests/mxfp4_ref.c tests/test_fp8_warp_cuda.cu tests/test_fp8_cuda.cu tests/test_weights_owned_cuda.cu tests/test_cuda_fmt_trap_cuda.cu tests/test_alloc_footprint_cuda.cu tests/test_cuda_init_failure.cu @command -v "$(GPUCC)" >/dev/null 2>&1 || { echo "$(GPUCC_NAME) not found" >&2; exit 1; } "$(GPUCC)" $(GPUFLAGS) backend_cuda.cu tests/test_backend_cuda.cu -o backend_cuda_test$(EXE) $(ANS_NVCC_LIBS) $(GPU_TEST_LIBS) ./backend_cuda_test$(EXE) @@ -1067,7 +1067,7 @@ gpu-compile: backend_cuda.o "$(GPUCC)" $(GPUFLAGS) tests/test_weights_owned_cuda.cu -o weights_owned_test$(EXE) $(ANS_NVCC_LIBS) $(GPU_TEST_LIBS) "$(GPUCC)" $(GPUFLAGS) tests/test_cuda_fmt_trap_cuda.cu -o cuda_fmt_trap_test$(EXE) $(ANS_NVCC_LIBS) $(GPU_TEST_LIBS) -cuda-bench: backend_cuda.cu backend_cuda.h backend_gpu_compat.h tests/bench_tensor_core.cu +cuda-bench: backend_cuda.cu backend_cuda.h fp8_format.h backend_gpu_compat.h tests/bench_tensor_core.cu @command -v "$(GPUCC)" >/dev/null 2>&1 || { echo "$(GPUCC_NAME) not found" >&2; exit 1; } "$(GPUCC)" $(GPUFLAGS) backend_cuda.cu tests/bench_tensor_core.cu -o backend_cuda_bench$(EXE) $(ANS_NVCC_LIBS) ./backend_cuda_bench$(EXE) @@ -1075,7 +1075,7 @@ cuda-bench: backend_cuda.cu backend_cuda.h backend_gpu_compat.h tests/bench_tens # fmt=8 kernel bench: old vs COLI_CUDA_F8_WARP kernels per decode path, census # expert shapes, JSON on stdout (kernel time only). The build is a separate # file target so tools/run_f8_bench.sh can keep stdout pure JSON. -fp8_bench$(EXE): backend_cuda.cu backend_cuda.h backend_gpu_compat.h tests/bench_fp8_cuda.cu +fp8_bench$(EXE): backend_cuda.cu backend_cuda.h fp8_format.h backend_gpu_compat.h tests/bench_fp8_cuda.cu @command -v "$(GPUCC)" >/dev/null 2>&1 || { echo "$(GPUCC_NAME) not found" >&2; exit 1; } "$(GPUCC)" $(GPUFLAGS) tests/bench_fp8_cuda.cu -o fp8_bench$(EXE) $(ANS_NVCC_LIBS) fp8-bench: fp8_bench$(EXE) @@ -1550,6 +1550,13 @@ tests/test_fp8_passthrough$(EXE): tests/test_fp8_passthrough.c quant.h fp8_forma tests/test_cuda_fmt_guard$(EXE): tests/test_cuda_fmt_guard.c backend_cuda.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) +# The fmt=8 LUT-gate state machine. Like test_cuda_fmt_guard it links no CUDA +# object and needs no toolchain: the two decisions live as pure predicates in +# backend_cuda.h and backend_cuda.cu calls them, so this pins the engine's own +# logic on a plain CPU build. +tests/test_cuda_lut_gate$(EXE): tests/test_cuda_lut_gate.c backend_cuda.h + $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) + tests/test_fp8_load$(EXE): tests/test_fp8_load.c colibri.c oracle.h st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h idot.h sample.h kv_persist.h telemetry.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) @@ -2082,6 +2089,13 @@ tests/test_olmoe_dot_i8_16$(EXE): tests/test_olmoe_dot_i8_16.c olmoe.c sse41_ker tests/test_olmoe_dot_i8_16_sse41$(EXE): tests/test_olmoe_dot_i8_16.c olmoe.c sse41_kernels.h st.h json.h compat.h sample.h tok.h tok_unicode.h tok_unicode_o200k.h omp_tune.h route_trace.h serve_codec.h $(CC) $(NOCUDA_CFLAGS) -msse4.1 -mno-avx2 -mno-fma $< -o $@ $(NOCUDA_LDFLAGS) +# layer_cuda_shard_kvb's format-allowlist refusal (see the test's file header). The +# function only exists under -DCOLI_CUDA, so on a default (CPU) build the binary is a +# loud SKIP; built with CUDA=1 it links backend_cuda.o ($(CUDA_OBJ)) and exercises the +# real guard -- no GPU needed, every probed path returns before any device context. +tests/test_shard_kvb_refuse$(EXE): tests/test_shard_kvb_refuse.c colibri.c st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h sample.h kv_persist.h telemetry.h $(CUDA_OBJ) + $(CC) $(CFLAGS) $< $(CUDA_OBJ) -o $@ $(LDFLAGS) + test-c: $(TEST_BINS) $(PYTHON) tools/run_tests.py $(TEST_BINS) diff --git a/c/backend_cuda.cu b/c/backend_cuda.cu index a0c892a52..b7930a0f7 100644 --- a/c/backend_cuda.cu +++ b/c/backend_cuda.cu @@ -1,7 +1,12 @@ #include "backend_cuda.h" +#include "fp8_format.h" /* FP8_BLOCK: the shared fmt=8 scale-block edge (see that header) */ #include "backend_gpu_compat.h" +static_assert(FP8_BLOCK == 128, "fmt=8 on-disk containers carry ceil(dim/128)-edged scale " + "grids (mint tool, docs/FORMATS.md); FP8_BLOCK is container format, not a " + "tunable -- an edit here is a format change"); + /* Optional fmt=8 decode candidate (COLI_CUDA_F8_WARP=2): cuda_fp8.h maps * __nv_cvt_fp8_to_halfraw to an sm_89+ cvt instruction, with a bit-manip * fallback below 890. CUDA-only; the HIP build keeps the LUT decode. */ @@ -287,13 +292,14 @@ __device__ static inline float mx4_weight_at(const uint8_t *q, int i) { * branch and the fall-through is a refusal. * * It used to be the other way round: int2 was the fall-through, so every format - * this function does not decode -- fmt=5 (int3-g64), fmt=6 (E8/IQ3), fmt=8 - * (fp8-e4m3), and anything added later -- was read as 2-bit values and returned - * numbers. Meanwhile the CPU functions doing the same job on the same tensor, - * qt_addrow and qt_matvec_rows (colibri.c), both exit(1) naming the function and - * the fmt. Two backends, identical unsupported input, one refusing and one - * fabricating: that asymmetry is the defect, independent of any particular - * format's arrival. + * this function does not decode -- fmt=5 (int3-g64), fmt=6 (E8/IQ3), and + * anything added later -- was read as 2-bit values and returned numbers. + * (fmt=8 was in that misread set too, then refused, until it gained its own + * explicit branch below for the absorb path.) Meanwhile the CPU functions + * doing the same job on the same tensor, qt_addrow and qt_matvec_rows + * (colibri.c), both exit(1) naming the function and the fmt. Two backends, + * identical unsupported input, one refusing and one fabricating: that + * asymmetry is the defect, independent of any particular format's arrival. * * WHY __trap() AND NOT A DIAGNOSTIC. This is device code inside a running * kernel; there is no stderr to name the tensor on and no way to unwind. __trap @@ -322,6 +328,15 @@ __device__ static float weight_at(const void *weights, int fmt, size_t row, int const uint8_t *base = static_cast(weights) + row; if (fmt == 0) return reinterpret_cast(base)[i]; if (fmt == 1) return static_cast(reinterpret_cast(base)[i]); + /* fmt=8 (fp8-e4m3): raw byte, same layout as fmt=1 (row_bytes(8,I)==I), decoded + * through the shared c_e4m3 LUT (same table quant_matmul's fmt==8 branch reads, + * uploaded once by coli_cuda_fp8_set_lut). Callers gate on the LUT being live + * before a fmt=8 tensor ever reaches this function (coli_cuda_tensor_upload + * refuses the upload otherwise), so the table is always populated here. Returns + * the decoded WEIGHT only, unscaled -- absorb_scale below applies the + * per-128x128-block scale, exactly like every other quantized fmt returns + * unscaled through this function. */ + if (fmt == 8) return c_e4m3[base[i]]; const uint8_t *q = base; if (fmt == 2 || fmt == 4) { /* fmt=4: same nibble layout */ uint8_t v = q[i >> 1]; @@ -335,13 +350,34 @@ __device__ static float weight_at(const void *weights, int fmt, size_t row, int return 0.0f; /* not reached: __trap() does not return */ } -/* Scale for output `row`, input element `k`. fmt=4 (grouped int4) stores ng - * scales per row at scales[row*ng + k/gs]; every other quantized format has - * one scale per row at scales[row]. Mirrors quant_matmul's fmt==4 branch so the +/* Scale for output `row`, input element `k`. Three layouts reach this: fmt=4 + * (grouped int4) stores ng scales per row at scales[row*ng + k/gs]; fmt=8 + * (fp8-e4m3) stores one scale per 128x128 BLOCK, block-row-major, and is + * handled by its own branch below; every OTHER quantized format has one scale + * per row at scales[row]. Mirrors quant_matmul's fmt==4 branch so the * attention absorb kernels apply per-group scales instead of the per-row * (fmt=2) semantic that crashed #298's g64 kv_b. */ __device__ static float absorb_scale(const float *wscale, int fmt, int gs, int ng, int row, int k) { if (!fmt) return 1.f; + if (fmt == 8) { + /* fp8-e4m3: one f32 scale per 128x128 BLOCK, block-row-major + * ([ceil(O/128), ceil(I/128)]), exactly quant_matmul's fmt==8 indexing + * (scl[i >> 7] on a scale row selected by o >> 7) and matmul_fp8's CPU + * reference (quant.h). `ng` here is coli_cuda_tensor_upload's t->ng, + * which for fmt=8 is set to ceil(I/128) specifically (not the fmt=4 + * group count) -- see the upload-time assignment there. `gs` is unused + * for fmt=8 (always 0, only fmt=4 sets it), so the block edge is the + * fixed FP8_BLOCK constant (fp8_format.h, shared with the CPU side), + * not a caller-supplied group size. Rounding note: the GEOMETRY here + * matches quant_matmul_f8w/matmul_fp8, but their fp8 accumulation + * convention (f32 partial per block, scale once per partial, double + * across blocks) is NOT carried into the absorb kernels -- they apply + * the scale per element into a float accumulator, matching their own + * fmt=4 arm's long-standing behavior; CPU-vs-CUDA absorb divergence + * is an accepted, documented class (#510). */ + int rowBlk = row / FP8_BLOCK, colBlk = k / FP8_BLOCK; + return wscale[(size_t)rowBlk * ng + colBlk]; + } if (fmt != 4) return wscale[row]; int g = k / gs; if (g >= ng) g = ng - 1; /* tail of the last (partial) group */ return wscale[(size_t)row * ng + g]; @@ -558,9 +594,9 @@ __global__ static void quant_matmul(float *y, const float *x, const void *weight * the ORIGINAL dense path, kept for COLI_CUDA_F8_WARP=0; the default * routes fmt=8 to quant_matmul_f8w instead (quant_matmul_launch). */ const uint8_t *wrow = static_cast(weights) + row; - const float *scl = scales + (size_t)(o >> 7) * (size_t)((I + 127) >> 7); + const float *scl = scales + (size_t)(o / FP8_BLOCK) * (size_t)((I + FP8_BLOCK - 1) / FP8_BLOCK); for (int i = threadIdx.x; i < I; i += blockDim.x) - sum += xs[i] * c_e4m3[wrow[i]] * scl[i >> 7]; + sum += xs[i] * c_e4m3[wrow[i]] * scl[i / FP8_BLOCK]; } else { for (int i = threadIdx.x; i < I; i += blockDim.x) sum += xs[i] * weight_at(weights, fmt, row, i); @@ -1272,10 +1308,15 @@ extern "C" int coli_cuda_init(const int *devices, int count) { } } if (g_nctx) { - int same = count == g_nctx; - for (int i = 0; same && i < count; i++) same = devices[i] == g_ctx[i].device; - if (!same) std::fprintf(stderr, "[CUDA] device list change requires shutdown first\n"); - return same; + /* Same decision as before, routed through the shared predicate in + * backend_cuda.h so a host-side test can pin it without nvcc; the + * return value is unchanged (1 for the same set, 0 otherwise). */ + int live[COLI_CUDA_MAX_DEVICES]; + for (int i = 0; i < g_nctx; i++) live[i] = g_ctx[i].device; + int d = coli_cuda_init_disposition(g_nctx, count, devices, live); + if (d == COLI_CUDA_INIT_REFUSE) + std::fprintf(stderr, "[CUDA] device list change requires shutdown first\n"); + return d == COLI_CUDA_INIT_ACCEPT; } for (int i = 0; i < count; i++) { int device = devices[i]; @@ -1356,6 +1397,22 @@ extern "C" void coli_cuda_shutdown(void) { ctx->group_desc=nullptr; ctx->group_desc_cap=0; } g_nctx = 0; + /* g_fp8_lut_ready is PROCESS-WIDE while the e4m3 table (c_e4m3, a + * __constant__ device symbol whose lifetime is the CUDA primary context, + * not this file's host-side DeviceContext structs) is PER-DEVICE. A later + * coli_cuda_init may select a device the previous span never published to; + * without this reset the upload gate (g_fp8_lut_ready, checked in + * coli_cuda_tensor_upload) would still be satisfied from the PREVIOUS + * boot and admit fmt=8 tensors whose kernels there decode against an + * unwritten (zero) table: silent all-zero weights, the exact + * fabricated-numbers failure mode the format gates exist to refuse. + * Reset so every boot must publish its own LUT (coli_cuda_fp8_set_lut) + * before any fmt=8 upload. Shutdown is the ONLY site that needs to clear + * the flag: coli_cuda_init refuses a re-init that names a different + * device set (it returns early while g_nctx is non-zero, leaving the + * existing contexts and their published table untouched), so the device + * set can only WIDEN by passing through here first. */ + g_fp8_lut_ready = 0; #ifdef COLI_ANS if(g_ans_sidecar){std::fclose(g_ans_sidecar);g_ans_sidecar=nullptr;} #if defined(__linux__) @@ -1444,7 +1501,9 @@ extern "C" int coli_cuda_tensor_upload(ColiCudaTensor **tensor, /* fmt=6 keeps its scales inside each 98-byte block, so it is the one * quantized format that legitimately arrives with scales == NULL. */ if (!rb || (fmt && fmt != 6 && !scales)) return 0; - if (fmt == 8 && !g_fp8_lut_ready) return 0; /* kernels would read a zero LUT */ + /* kernels would read a zero LUT; shared predicate, pinned by + * tests/test_cuda_lut_gate.c without a CUDA toolchain */ + if (!coli_cuda_fp8_gate_admits(fmt, g_fp8_lut_ready)) return 0; ColiCudaTensor *t = static_cast(std::calloc(1, sizeof(*t))); if (!t) return 0; t->fmt = fmt; t->I = I; t->O = O; t->device = device; t->weight_bytes = rb * (size_t)O; @@ -1455,9 +1514,9 @@ extern "C" int coli_cuda_tensor_upload(ColiCudaTensor **tensor, t->ng = (I + 31) / 32; t->scale_count = (size_t)O * t->ng; } - if (fmt == 8) { /* per-128x128-block scales: [ceil(O/128), ceil(I/128)] */ - t->ng = (I + 127) / 128; - t->scale_count = (size_t)((O + 127) / 128) * (size_t)t->ng; + if (fmt == 8) { /* per-block scales: [ceil(O/FP8_BLOCK), ceil(I/FP8_BLOCK)] (fp8_format.h) */ + t->ng = (int)fp8_nblk(I); + t->scale_count = (size_t)fp8_nblk(O) * (size_t)t->ng; } if (!cuda_ok(cudaMalloc(&t->weights, t->weight_bytes), "tensor allocation")) { coli_cuda_tensor_free(t); @@ -2278,18 +2337,20 @@ extern "C" const float *coli_cuda_expert_group_take(int device) { /* The absorb kernels decode `w` through weight_at + absorb_scale, which know - * per-row and fmt=4 group scales only. Refuse anything else (fmt=5/6/8) rather - * than mis-decode it — the caller keeps its CPU attention path. (`proj` - * tensors are exempt: they run through quant_matmul, which dispatches every - * format it uploads.) A dedicated block-scale absorb for fmt=8 is follow-up - * work, same shape as routing fmt=4 through the grouped kernels was. + * per-row scales, fmt=4 group scales, and fmt=8 per-128x128-block scales. + * Refuse anything else (fmt=5/6/7) rather than mis-decode it — the caller + * keeps its CPU attention path. (`proj` tensors are exempt: they run through + * quant_matmul, which dispatches every format it uploads.) fmt=8 support + * funnels through this one predicate for all the absorb host wrappers below, + * so none of them needed a separate change. * * The admissible set is weight_at's own, taken from the shared predicate rather - * than restated as `fmt <= 4`: this gate and weight_at's device-side backstop - * must not be able to drift apart, and the old inequality also admitted - * NEGATIVE fmt values, which weight_at would then have fallen through on. Same - * truth table for every fmt a container can actually carry (0..8), so no - * existing container changes behaviour here. */ + * than restated as an inequality: this gate and weight_at's device-side + * backstop must not be able to drift apart, and the old `fmt <= 4` also + * admitted NEGATIVE fmt values, which weight_at would then have fallen through + * on. A fmt=8 tensor implies a live e4m3 LUT (upload refuses it otherwise -- + * see the predicate's caveat note in backend_cuda.h), so no extra gate is + * needed here. */ static int absorb_fmt_ok(const ColiCudaTensor *w){ return w && coli_cuda_weight_at_supported(w->fmt); } diff --git a/c/backend_cuda.h b/c/backend_cuda.h index 2579561d2..e58dbf95f 100644 --- a/c/backend_cuda.h +++ b/c/backend_cuda.h @@ -23,7 +23,9 @@ extern "C" { /* Weight formats the generic per-element device decoder (weight_at, * backend_cuda.cu) can actually decode: f32, int8-row, int4 nibbles (fmt=2 and - * the grouped fmt=4, same packing), and int2. Nothing else. + * the grouped fmt=4, same packing), int2, and fmt=8 (fp8-e4m3 raw bytes, + * decoded through the c_e4m3 LUT -- absorb-path support; absorb_scale supplies + * its per-128x128-block scale). Nothing else. * * WHY THIS IS A PREDICATE AND NOT A COMMENT. weight_at used to END in the int2 * decode as an unguarded fall-through, so ANY other format handed to it -- a @@ -42,12 +44,60 @@ extern "C" { * same arrangement colibri.c uses for metal_fused_fmt_ok. * * NOT a statement about which formats the CUDA BACKEND supports: quant_matmul - * has its own explicit branches for fmt=6 (E8/IQ3), fmt=7 (MXFP4) and fmt=8 - * (fp8-e4m3) that never route through weight_at. This predicate is scoped to - * weight_at's own dispatch, which is what the absorb and grouped-expert kernels - * decode through. */ + * has its own explicit branches for fmt=6 (E8/IQ3) and fmt=7 (MXFP4) that + * never route through weight_at (and its own fmt=8 branch for the dense path + * -- weight_at's fmt=8 branch serves the absorb kernels, which share the same + * c_e4m3 LUT). This predicate is scoped to weight_at's own dispatch, which is + * what the absorb and grouped-expert kernels decode through. + * + * fmt=8 CAVEAT, stated because the truth table alone cannot carry it: a fmt=8 + * decode additionally requires the e4m3 LUT to have been published to the + * configured devices (coli_cuda_fp8_set_lut). The exact mechanism, so the + * claim cannot outrun it: coli_cuda_fp8_set_lut copies the table into every + * context live AT CALL TIME and sets a process-wide flag; the flag gates + * fmt=8 uploads (coli_cuda_tensor_upload refuses until it is set). + * coli_cuda_shutdown clears the flag; coli_cuda_init never writes it. That is + * enough because init will not rebuild contexts underneath a live set: a + * re-init naming the same device set returns success and leaves the contexts, + * and the table published to them, untouched, while one naming a different + * device set is refused before any context is rebuilt. So the device set + * cannot widen past what the last publish covered without going through + * shutdown, and no fmt=8 ColiCudaTensor can reach a kernel whose device has + * an unwritten table. This predicate + * deliberately does not restate that gate: it answers "does weight_at have a + * decode branch for this fmt", which is the question the launch-site gates + * and the device-side __trap() backstop share. */ static inline int coli_cuda_weight_at_supported(int fmt) { - return fmt == 0 || fmt == 1 || fmt == 2 || fmt == 3 || fmt == 4; + return fmt == 0 || fmt == 1 || fmt == 2 || fmt == 3 || fmt == 4 || fmt == 8; +} + +/* The two decisions the fmt=8 LUT gate rests on, as pure predicates. They live + * here rather than inline in backend_cuda.cu so a host-side test can pin them + * with no CUDA toolchain and no GPU (tests/test_cuda_lut_gate.c). backend_cuda.cu + * calls BOTH at the real decision sites, so the test pins the engine's own + * logic rather than a second copy that could drift from it -- which is the + * failure this factoring exists to prevent, the gate having no CI reach + * otherwise. */ + +/* Does the upload gate admit this tensor? Only fmt=8 needs the published + * table; every other format decodes without one. */ +static inline int coli_cuda_fp8_gate_admits(int fmt, int lut_ready) { + return fmt != 8 || lut_ready != 0; +} + +/* What coli_cuda_init must do with a request while a device set may be live. + * BUILD: nothing is live, build the contexts. ACCEPT: the same set is already + * live -- return success and touch nothing, so the table published to those + * contexts stays valid. REFUSE: a different set is live -- refuse before + * rebuilding anything, so the set cannot widen past the last publish. */ +enum { COLI_CUDA_INIT_BUILD = 0, COLI_CUDA_INIT_ACCEPT = 1, COLI_CUDA_INIT_REFUSE = -1 }; +static inline int coli_cuda_init_disposition(int nctx, int count, + const int *want, const int *live) { + int i; + if (nctx <= 0) return COLI_CUDA_INIT_BUILD; + if (count != nctx) return COLI_CUDA_INIT_REFUSE; + for (i = 0; i < count; i++) if (want[i] != live[i]) return COLI_CUDA_INIT_REFUSE; + return COLI_CUDA_INIT_ACCEPT; } /* Opaque, persistent device copy of one resident quantized tensor. */ diff --git a/c/colibri.c b/c/colibri.c index 8ffd06bde..9d91c4611 100644 --- a/c/colibri.c +++ b/c/colibri.c @@ -2278,10 +2278,9 @@ static void layer_cuda_shard_kvb(Layer *l,int H,int Q,int V){ * dereferenced it -- silent, unnamed, and one refactor away from a misread. * Refuse BY NAME instead, BEFORE any pointer/stride use, and say what happens * instead: the un-sharded kv_b stays whole on its layer home device, where fmt=8 - * kv_b decode runs the CPU absorb path (qt_addrow/qt_matvec_rows' fmt=8 - * branches below; the CUDA absorb kernels refuse fmt=8 via absorb_fmt_ok's - * fmt 0..4 allowlist) -- COLI_CUDA_ATTN_SHARD is a no-op for it. Same - * "refuse rather than misread" discipline as qt_addrow/ + * kv_b decode runs the absorb path (qt_addrow/qt_matvec_rows' fmt=8 branches, or + * the CUDA absorb kernels via absorb_fmt_ok) -- COLI_CUDA_ATTN_SHARD is a no-op + * for it. Same "refuse rather than misread" discipline as qt_addrow/ * qt_matvec_rows' guards; notice only (no exit): sharding is an opt-in * optimization and skipping it is the correct, working behavior. Bounded once * per process per fmt, never per layer (metal_fmt_gate_notice, the precedent diff --git a/c/fp8_format.h b/c/fp8_format.h index 04610efef..d10ea461e 100644 --- a/c/fp8_format.h +++ b/c/fp8_format.h @@ -1,20 +1,25 @@ /* fmt=8 (fp8-e4m3-b128) block geometry -- the single definition site for the - * 128x128 scale-block edge. Every consumer reaches these symbols through - * quant.h, which includes this header where the definitions used to live: - * matmul_fp8 and the e4m3 dequant plumbing (quant.h), the qt_addrow / - * qt_matvec_rows / qt_from_disk paths (colibri.c), and the native-FP8 - * reduction (qwen38.c / qwen38_core.h). Kept deliberately tiny (no LUTs, no - * functions with OpenMP pragmas, no intrinsics) so a translation unit that - * cannot include quant.h itself can still name the block edge instead of - * restating 128 / >>7 as literals. + * 128x128 scale-block edge, shared by the CPU side (quant.h: matmul_fp8, + * e4m3 dequant plumbing; colibri.c: qt_addrow/qt_matvec_rows, qt_from_disk) + * and the CUDA backend's three converted sites (backend_cuda.cu: + * absorb_scale, quant_matmul's fmt=8 branch, the upload-time ng/scale_count + * computation); the f8-warp and f8-group kernels there keep their own + * 128 / >>7 literals, pinned against drift by backend_cuda.cu's + * static_assert(FP8_BLOCK == 128). Before this header the two backends + * agreed by IDENTICAL LITERALS restated in each file, so an edit to one side + * could not break the other side's build -- only a full-scale CPU-vs-CUDA + * parity run would have noticed. Kept deliberately tiny (no LUTs, no + * functions with OpenMP pragmas, no intrinsics) so the CUDA translation unit + * can include it without dragging in quant.h. * * The in-repo FP8 tests are NOT independent of this constant: the reference * decoders in test_fp8_passthrough.c, test_fp8_load.c, test_fp8_e2e_loader.c, - * test_qwen38_native_weights.c and test_qt_addrow.c all index with - * FP8_BLOCK/fp8_nblk via quant.h, so they move in lockstep with an edit here - * (only test_backend_metal.mm's ref_fp8_nblk keeps its own literal, on purpose). - * An edit to FP8_BLOCK is a format change, not a tunable -- those tests - * cannot catch a wrong edit on their own. */ + * test_qwen38_native_weights.c, test_qt_addrow.c and test_shard_kvb_refuse.c + * all index with FP8_BLOCK/fp8_nblk via quant.h, so they move in lockstep + * with an edit here (test_backend_metal.mm's ref_fp8_nblk and + * test_backend_cuda.cu's fmt=8 reference keep their own literals, on + * purpose). An edit to FP8_BLOCK is a format change, not a tunable -- those + * tests cannot catch a wrong edit on their own. */ #ifndef COLI_FP8_FORMAT_H #define COLI_FP8_FORMAT_H @@ -22,9 +27,9 @@ #define FP8_BLOCK 128 -/* Blocks covering n elements: ceil(n/FP8_BLOCK). Plain C on purpose -- - * nothing in this header may need decoration or a special front end, so it - * stays includable from any translation unit. */ +/* Blocks covering n elements: ceil(n/FP8_BLOCK). Host-side helper (device + * code uses the FP8_BLOCK macro arithmetic directly -- this is not decorated + * for device compilation on purpose, to keep the header plain C). */ static inline int64_t fp8_nblk(int n){ return ((int64_t)n + FP8_BLOCK - 1) / FP8_BLOCK; } #endif /* COLI_FP8_FORMAT_H */ diff --git a/c/tests/test_backend_cuda.cu b/c/tests/test_backend_cuda.cu index 049ccf0a0..1603048c7 100644 --- a/c/tests/test_backend_cuda.cu +++ b/c/tests/test_backend_cuda.cu @@ -223,6 +223,106 @@ static int test_fmt6(int dev) { return 1; } +/* ---- fmt=8 (fp8-e4m3) absorb decode ----------------------------------- + * Exercises weight_at's new fmt=8 branch and absorb_scale's new per-128x128- + * block branch (this PR pair) through the REAL attention_absorb kernel and + * absorb_fmt_ok gate, against a CPU reference built the same way the fmt=0 + * absorb block in main() (below) is: independent score/softmax/context + * accumulation, only the weight lookup itself changes to an e4m3 block-scale + * dequant. Dims are chosen so BOTH the row-block and column-block axes carry + * a partial tail block (O=H*(Q+V)=160 -> nblkO=2, rows 128..159 partial; + * K=140 -> nblkI=2, cols 128..139 partial) -- the exact geometry the new + * branches must index correctly (blkO=row>>7, blkI=k>>7, scale index + * blkO*nblkI+blkI). The e4m3 reference decoder is arithmetic (sign/exp/mant, + * OCP E4M3-FN policy), not the engine's c_e4m3 LUT, so it cross-checks + * coli_cuda_fp8_set_lut's uploaded table rather than assuming it -- same + * independence discipline as t8_e4m3_ref's sibling in tests/test_fp8_cuda.cu. */ +static float t8_e4m3_ref(uint8_t b) { + int s = b >> 7, e = (b >> 3) & 15, m = b & 7; + if (e == 15 && m == 7) return NAN; /* E4M3-FN: only NaN, no inf */ + float v = e ? ldexpf(1.f + m/8.f, e-7) : ldexpf(m/8.f, -6); + return s ? -v : v; +} +static uint32_t t8_rng_state = 0xC001D00Du; +static uint8_t t8_rnd_byte(void) { + t8_rng_state ^= t8_rng_state<<13; t8_rng_state ^= t8_rng_state>>17; t8_rng_state ^= t8_rng_state<<5; + uint8_t b = (uint8_t)(t8_rng_state & 0xFF); + if ((b & 0x7F) == 0x7F) b &= (uint8_t)~1; /* avoid the two NaN byte patterns */ + return b; +} +static float t8_dequant(const uint8_t *q, const float *scale, int I, int row, int col, int nblkI) { + int blkO = row >> 7, blkI = col >> 7; + return t8_e4m3_ref(q[(size_t)row*I + col]) * scale[(size_t)blkO*nblkI + blkI]; +} + +static int test_fmt8_absorb(int dev) { + float lut[256]; for (int i = 0; i < 256; i++) lut[i] = t8_e4m3_ref((uint8_t)i); + if (!coli_cuda_fp8_set_lut(lut)) { std::fprintf(stderr,"fmt=8 absorb: set_lut failed\n"); return 0; } + + const int H = 2, Q = 40, V = 40, R = 2, K = 140, T = 3, O = H*(Q+V); + const int nblkO = (O+127)/128, nblkI = (K+127)/128, nblk = nblkO*nblkI; + uint8_t *w = (uint8_t*)std::malloc((size_t)O*K); + for (size_t i = 0; i < (size_t)O*K; i++) w[i] = t8_rnd_byte(); + float *wscale = (float*)std::malloc((size_t)nblk*sizeof(float)); + for (int i = 0; i < nblk; i++) wscale[i] = 0.01f + 0.002f*(float)i; + + /* Refusal must have held BEFORE this test ever ran (fmt=8 was invisible to + * absorb_fmt_ok on unpatched main) -- upload + launch below is the positive + * side of the same predicate this PR widened. */ + ColiCudaTensor *wt = nullptr; + if (!coli_cuda_tensor_upload(&wt, w, wscale, 8, K, O, dev)) { + std::fprintf(stderr,"fmt=8 absorb: weight upload rejected\n"); return 0; + } + + float *q = (float*)std::malloc((size_t)H*(Q+R)*sizeof(float)); + float *latent = (float*)std::malloc((size_t)T*K*sizeof(float)); + float *rope = (float*)std::malloc((size_t)T*R*sizeof(float)); + for (int i = 0; i < H*(Q+R); i++) q[i] = std::sin((float)(i+1)*0.037f); + for (int i = 0; i < T*K; i++) latent[i] = std::sin((float)(i+1)*0.019f)*0.5f; + for (int i = 0; i < T*R; i++) rope[i] = std::cos((float)(i+1)*0.041f)*0.3f; + float *ctx = (float*)std::malloc((size_t)H*V*sizeof(float)); + float scale = 1.f/std::sqrt((float)K); + + if (!coli_cuda_attention_absorb(wt, ctx, q, latent, rope, H, Q, R, V, K, T, scale)) { + std::fprintf(stderr,"fmt=8 absorb: kernel launch rejected (absorb_fmt_ok gate?)\n"); return 0; + } + + int bad = 0; + for (int h = 0; h < H; h++) { + int rbase = h*(Q+V); + float qa[512]; /* K<=512, this attention_absorb's own documented bound */ + for (int k = 0; k < K; k++) { + double a = 0; + for (int d = 0; d < Q; d++) a += (double)q[h*(Q+R)+d]*t8_dequant(w,wscale,K,rbase+d,k,nblkI); + qa[k] = (float)a; + } + float scores[T]; + for (int t = 0; t < T; t++) { + double a = 0; for (int k = 0; k < K; k++) a += (double)qa[k]*latent[t*K+k]; + for (int d = 0; d < R; d++) a += (double)q[h*(Q+R)+Q+d]*rope[t*R+d]; + scores[t] = (float)a*scale; + } + float mx = scores[0]; for (int t = 1; t < T; t++) mx = scores[t]>mx?scores[t]:mx; + float z = 0; for (int t = 0; t < T; t++) { scores[t] = std::exp(scores[t]-mx); z += scores[t]; } + for (int t = 0; t < T; t++) scores[t] /= z; + float cl[512]; + for (int k = 0; k < K; k++) { double a=0; for (int t=0;t1e-4f ? std::fabs(got-want)/std::fabs(want) : std::fabs(got-want); + if (rel > 1e-3f) { + std::fprintf(stderr,"fmt=8 absorb mismatch h=%d v=%d got=%.6f want=%.6f rel=%.4g\n",h,v,got,want,rel); + bad++; + } + } + } + coli_cuda_tensor_free(wt); + std::free(w); std::free(wscale); std::free(q); std::free(latent); std::free(rope); std::free(ctx); + return bad == 0; +} + int main(int argc, char **argv) { int devices[COLI_CUDA_MAX_DEVICES], ndev = argc > 1 ? argc - 1 : 1; if (ndev > COLI_CUDA_MAX_DEVICES) return 2; @@ -455,6 +555,20 @@ int main(int argc, char **argv) { coli_cuda_stats(-1, &count, &bytes); if (count || bytes) { std::fprintf(stderr,"fmt=6 leaked tensors\n"); return 1; } + /* fmt=8 absorb: same self-contained lifecycle discipline as fmt=6 above, + * BOTH halves. The byte half is load-bearing history: an earlier vintage of + * this feature found coli_cuda_tensor_free subtracting per-row scale bytes + * for a per-block-scaled fmt=8 tensor (upload charged scale_count = + * ceil(O/128)*ng, free subtracted O*ng), so the `tensor_bytes >= bytes` + * guard silently declined the subtraction and the diagnostic VRAM counter + * stuck non-zero forever after freeing ANY fmt=8 tensor. free's accounting + * now mirrors upload's charge expression exactly (see the comment in + * coli_cuda_tensor_free), and this is the assertion that keeps the two + * from drifting apart again for a tracked fmt=8 tensor. */ + if (!test_fmt8_absorb(d0)) return 1; + coli_cuda_stats(-1, &count, &bytes); + if (count || bytes) { std::fprintf(stderr,"fmt=8 absorb leaked tensors\n"); return 1; } + coli_cuda_shutdown(); std::printf("cuda backend: q8/q4/q2/f32/e8 correctness ok on %d device(s)\n", ndev); return 0; diff --git a/c/tests/test_cuda_fmt_guard.c b/c/tests/test_cuda_fmt_guard.c index 056d68535..85e5fee60 100644 --- a/c/tests/test_cuda_fmt_guard.c +++ b/c/tests/test_cuda_fmt_guard.c @@ -36,21 +36,27 @@ static int fails = 0; #define CHECK(c) do{ if(!(c)){ printf("FAIL %s:%d: %s\n", __FILE__, __LINE__, #c); fails++; } }while(0) int main(void) { - /* Decodable by weight_at's explicit branches. */ + /* Decodable by weight_at's explicit branches. fmt=8 (fp8-e4m3-b128) joined + * this set when the absorb path gained its block-scale decode (weight_at's + * fmt==8 branch + absorb_scale's per-128x128-block branch); its decode + * additionally needs the e4m3 LUT live on the device, which is an + * UPLOAD-time gate (coli_cuda_tensor_upload refuses fmt=8 until + * coli_cuda_fp8_set_lut has run), deliberately not part of this truth + * table -- see the predicate's own caveat note in backend_cuda.h. */ CHECK(coli_cuda_weight_at_supported(0)); /* f32 */ CHECK(coli_cuda_weight_at_supported(1)); /* int8-row */ CHECK(coli_cuda_weight_at_supported(2)); /* int4 nibbles */ CHECK(coli_cuda_weight_at_supported(3)); /* int2 */ CHECK(coli_cuda_weight_at_supported(4)); /* grouped int4 */ + CHECK(coli_cuda_weight_at_supported(8)); /* fp8-e4m3-b128 */ /* NOT decodable. Each of these has a real in-tree meaning, and each used to * be read as int2 by the fall-through. fmt=5 and fmt=6 are the ones that - * could already reach a CUDA tensor; fmt=8 is the one this branch is about; - * fmt=7 (MXFP4) has its own quant_matmul branch and never routes here. */ + * could already reach a CUDA tensor; fmt=7 (MXFP4) has its own quant_matmul + * branch and never routes here. */ CHECK(!coli_cuda_weight_at_supported(5)); /* int3-g64 */ CHECK(!coli_cuda_weight_at_supported(6)); /* E8/IQ3 */ CHECK(!coli_cuda_weight_at_supported(7)); /* MXFP4 */ - CHECK(!coli_cuda_weight_at_supported(8)); /* fp8-e4m3-b128 */ /* Out of range in both directions. The negative cases are the ones the * previous `fmt <= 4` host gate admitted: a descriptor whose fmt field is @@ -62,10 +68,10 @@ int main(void) { CHECK(!coli_cuda_weight_at_supported(100)); /* PRIVATE ORDINAL BLOCK base */ CHECK(!coli_cuda_weight_at_supported(1 << 30)); - /* The set is exactly {0,1,2,3,4} and nothing else in a wide sweep -- so a + /* The set is exactly {0,1,2,3,4,8} and nothing else in a wide sweep -- so a * later edit that widens the predicate has to change this line too. */ for (int fmt = -2048; fmt <= 2048; fmt++) { - int expect = (fmt >= 0 && fmt <= 4); + int expect = (fmt >= 0 && fmt <= 4) || fmt == 8; if (!!coli_cuda_weight_at_supported(fmt) != expect) { printf("FAIL fmt=%d: supported=%d, expected %d\n", fmt, coli_cuda_weight_at_supported(fmt), expect); diff --git a/c/tests/test_cuda_fmt_trap_cuda.cu b/c/tests/test_cuda_fmt_trap_cuda.cu index 07b2372f1..78e22691e 100644 --- a/c/tests/test_cuda_fmt_trap_cuda.cu +++ b/c/tests/test_cuda_fmt_trap_cuda.cu @@ -31,7 +31,7 @@ * WHAT MAKES THIS TEST BITE RATHER THAN MERELY PASS. Three controls, because an * exit code that says "the child failed" is worthless if the child fails no * matter what: - * 1. IN-PROCESS control: every supported fmt (0,1,2,3,4) is launched in the + * 1. IN-PROCESS control: every supported fmt (0,1,2,3,4,8) is launched in the * parent and must complete cleanly. If weight_at trapped on those, the * trap would be firing on valid input and this test would fail here. * 2. CHILD-HARNESS control: the probe list always contains SUPPORTED formats @@ -56,13 +56,24 @@ * landing in either order: whatever the predicate promises, the silicon must * deliver, and whatever it refuses, the silicon must refuse. * - * The fmt set probed covers both sides of today's truth table (0,3 supported; - * 5,6,7,8 unsupported container-carriable formats) plus the out-of-range values - * a corrupt descriptor could present (-1, 9, 1<<30) -- the ones the previous + * The fmt set probed covers both sides of today's truth table (0,3,8 supported + * -- 8 since the fp8-e4m3 absorb decode widened the predicate; 5,6,7 + * unsupported container-carriable formats) plus the out-of-range values a + * corrupt descriptor could present (-1, 9, 1<<30) -- the ones the previous * `fmt <= 4` host gate admitted. * + * fmt=8 gets one extra obligation on top of the derived DECODED verdict: its + * decoded VALUE is checked against an independent arithmetic e4m3 reference + * (a6_e4m3_ref below), because a fresh child process starts with a + * zero-initialized c_e4m3 table -- a decode that "returns a number" from a + * zero LUT is exactly the fabricated-numbers failure mode this file exists to + * refuse. The child publishes the LUT through the real + * coli_cuda_init/coli_cuda_fp8_set_lut path first, the same way any real + * fmt=8 caller must before upload. + * * make -C c cuda-test CUDA_ARCH=native (runs this among the CUDA tests) */ +#include #include #include #include @@ -95,6 +106,8 @@ #define A6_EXIT_TRAPPED 42 /* the launch aborted -- the refusal fired */ #define A6_EXIT_DECODED 43 /* the kernel returned a value -- no refusal */ #define A6_EXIT_HARNESS 44 /* could not run the probe at all */ +#define A6_EXIT_MISMATCH 45 /* decoded without trapping, but not to the * + * value the independent reference expects */ #define A6_PROBE_FLAG "--a6-probe" @@ -103,6 +116,38 @@ * visible number rather than something that could be mistaken for "no output". */ #define A6_FILL 0xA5 +/* Independent arithmetic e4m3 decoder (sign/exp/mantissa, OCP E4M3-FN: no + * infinities, only 0x7F/0xFF are NaN) -- same construction as t8_e4m3_ref in + * tests/test_backend_cuda.cu and e4m3_ref in tests/test_fp8_cuda.cu. Scope + * of the check, stated exactly: the device table is PUBLISHED FROM this + * reference (a6_publish_lut), so the comparison proves the table is nonzero + * and that weight_at's c_e4m3[base[i]] indexing is correct -- it cannot + * catch a wrong reference (a wrong table built from it would cancel), and + * the engine's own E4M3_LUT (quant.h) is not exercised by this harness at + * all. A6_FILL (0xA5) is not a NaN byte pattern. */ +static float a6_e4m3_ref(uint8_t b) { + int s = b >> 7, e = (b >> 3) & 15, m = b & 7; + if (e == 15 && m == 7) return NAN; /* E4M3-FN: only NaN, no inf */ + float v = e ? ldexpf(1.f + m/8.f, e-7) : ldexpf(m/8.f, -6); + return s ? -v : v; +} + +/* Publish the e4m3 LUT through the engine's own path, so weight_at's fmt=8 + * branch reads a real table instead of context-fresh zeros. Harmless for every + * other fmt (only the fmt==8 branch reads c_e4m3). coli_cuda_fp8_set_lut walks + * the engine's context table (g_nctx/g_ctx), which nothing else in this file + * populates -- the rest of the file talks to the device directly via the raw + * CUDA runtime API, deliberately, to stay independent of the engine's + * device-selection plumbing -- so coli_cuda_init(device 0) is called for this + * one dependency only. Returns 0 on harness failure. */ +static int a6_publish_lut(void) { + int dev0 = 0; + if (!coli_cuda_init(&dev0, 1)) return 0; + float lut[256]; + for (int i = 0; i < 256; i++) lut[i] = a6_e4m3_ref((uint8_t)i); + return coli_cuda_fp8_set_lut(lut); +} + /* The deliberate bad launch. weight_at is file-static device code; this is the * only caller in this TU, and it passes fmt straight through as a runtime * argument so nvcc cannot constant-fold the dispatch away. */ @@ -116,6 +161,16 @@ static int a6_probe_child(int fmt) { void *w = NULL; float *out = NULL; + /* Fresh execv'd process, fresh CUDA context: c_e4m3 starts zero-initialized + * here regardless of what the parent published. Published unconditionally + * rather than gated on fmt==8, so this child's setup mirrors a real + * caller's -- the engine publishes the LUT once at boot for every process + * that might touch fmt=8, not per tensor. */ + if (!a6_publish_lut()) { + printf(" [child fmt=%d] coli_cuda_init/coli_cuda_fp8_set_lut failed\n", fmt); + return A6_EXIT_HARNESS; + } + if (cudaMalloc(&w, 256) != cudaSuccess) { printf(" [child fmt=%d] cudaMalloc(weights) failed\n", fmt); return A6_EXIT_HARNESS; @@ -155,6 +210,24 @@ static int a6_probe_child(int fmt) { return A6_EXIT_TRAPPED; } printf(" [child fmt=%d] weight_at RETURNED %g\n", fmt, (double)v); + + /* fmt=8's extra obligation (see the file header): the decoded value must + * match the independent e4m3 reference. Without this, a zero-LUT decode + * (or any wrong LUT/indexing) would still count as "DECODED" -- a wrong + * number is not meaningfully different from a fabricated one, which is + * the exact failure mode this file's __trap() backstop exists to refuse. + * The other DECODED fmts in the probe list (0, 3) are already pinned + * bit-exact elsewhere (inprocess_supported_control here, the dense-matmul + * oracle in tests/test_backend_cuda.cu). */ + if (fmt == 8) { + float want = a6_e4m3_ref((uint8_t)A6_FILL); + if (v != want) { + printf(" [child fmt=%d] MISMATCH: weight_at(A6_FILL=0x%02X) = %g, " + "independent e4m3 reference says %g\n", fmt, A6_FILL, + (double)v, (double)want); + return A6_EXIT_MISMATCH; + } + } return A6_EXIT_DECODED; } @@ -210,7 +283,11 @@ static void expect_child(const char *self, int fmt, int want, const char *why) { printf("ok fmt=%-6d %s (child exit %d)\n", fmt, why, got); return; } - if (got == A6_EXIT_DECODED && want == A6_EXIT_TRAPPED) { + if (got == A6_EXIT_MISMATCH && want == A6_EXIT_DECODED) { + printf("FAIL fmt=%-6d %s -- decoded without trapping, but not to the " + "value the independent e4m3 reference expects (see child output " + "above)\n", fmt, why); + } else if (got == A6_EXIT_DECODED && want == A6_EXIT_TRAPPED) { printf("FAIL fmt=%-6d weight_at DECODED a format the host predicate " "refuses -- the device-side __trap() is not firing (dispatch " "disagrees with coli_cuda_weight_at_supported)\n", fmt); @@ -247,7 +324,14 @@ static void inprocess_supported_control(void) { return; } (void)cudaGetLastError(); - for (int fmt = 0; fmt <= 4; fmt++) { + /* The full predicate-true set, fmt=8 included -- its LUT is published by + * main() before this control runs (a6_publish_lut), so its in-process + * decode reads a real table, and its value is checked against the + * independent reference right here (the child probes re-check it in a + * fresh process, where the LUT must be re-published). */ + static const int supported[] = {0, 1, 2, 3, 4, 8}; + for (size_t fi = 0; fi < sizeof supported / sizeof supported[0]; fi++) { + int fmt = supported[fi]; if (!coli_cuda_weight_at_supported(fmt)) { /* keeps the two in step */ printf("FAIL fmt=%d is in this loop but the host predicate rejects it\n", fmt); fails++; @@ -267,6 +351,12 @@ static void inprocess_supported_control(void) { fails++; return; /* the context is poisoned; nothing after this is valid */ } + if (fmt == 8 && v != a6_e4m3_ref((uint8_t)A6_FILL)) { + printf("FAIL fmt=8 decoded in-process to %g, independent e4m3 " + "reference says %g (wrong LUT or wrong indexing)\n", + (double)v, (double)a6_e4m3_ref((uint8_t)A6_FILL)); + fails++; + } printf("ok fmt=%-6d supported: decoded in-process, weight_at = %g\n", fmt, (double)v); } @@ -288,6 +378,13 @@ int main(int argc, char **argv) { return 1; } + if (!a6_publish_lut()) { + printf("cuda fmt trap test: could not publish the e4m3 LUT " + "(coli_cuda_init/coli_cuda_fp8_set_lut failed) -- the fmt=8 " + "control cannot run\n"); + return 1; + } + printf("-- control: supported formats decode (in-process)\n"); inprocess_supported_control(); diff --git a/c/tests/test_cuda_lut_gate.c b/c/tests/test_cuda_lut_gate.c new file mode 100644 index 000000000..ab6a1d5c3 --- /dev/null +++ b/c/tests/test_cuda_lut_gate.c @@ -0,0 +1,142 @@ +/* fmt=8 LUT-gate state machine, pinned on the host with no CUDA toolchain. + * + * Why this exists: the gate that stops a fmt=8 tensor reaching a kernel whose + * device has an unwritten e4m3 table is implemented in backend_cuda.cu, which + * compiles only under nvcc/hipcc. Every test that exercised it therefore ran + * nowhere in CPU CI, and a rebase was able to change the surrounding + * init semantics while three sites in the tree went on asserting the old ones. + * The two decisions the gate rests on are pure predicates in backend_cuda.h; + * backend_cuda.cu calls them at the real sites, so what is pinned below is the + * engine's own logic, not a restatement of it. + * + * What this does NOT pin: that backend_cuda.cu calls the predicates in the + * right order, or at all. That needs a CUDA toolchain (tests/test_fp8_cuda.cu + * Phase 3 covers it there). What it does pin is the decision table itself and + * the one invariant the whole gate rests on: + * + * no live context set => the LUT flag is clear + * + * which holds because shutdown is the only writer that clears the flag and it + * clears g_nctx in the same breath, and because init refuses to rebuild + * contexts underneath a live set. + */ +#include +#include "../backend_cuda.h" + +static int fails; + +static void ck(int cond, const char *what) { + if (!cond) { printf("FAIL %s\n", what); fails++; } +} + +/* A model of the three public transitions, written in terms of the SAME + * predicates the engine calls. `nctx`/`ready` are the engine's two globals. */ +typedef struct { int nctx; int ready; int dev[8]; } Gate; + +static int g_init(Gate *g, const int *want, int count) { + int d = coli_cuda_init_disposition(g->nctx, count, want, g->dev); + if (d == COLI_CUDA_INIT_REFUSE) return 0; /* touches nothing */ + if (d == COLI_CUDA_INIT_ACCEPT) return 1; /* touches nothing */ + for (int i = 0; i < count; i++) g->dev[i] = want[i]; + g->nctx = count; /* flag NOT written here */ + return 1; +} +static void g_set_lut(Gate *g) { if (g->nctx > 0) g->ready = 1; } +static void g_shutdown(Gate *g) { g->nctx = 0; g->ready = 0; } +static int g_upload(Gate *g, int fmt) { + return g->nctx > 0 && coli_cuda_fp8_gate_admits(fmt, g->ready); +} + +int main(void) { + /* --- the disposition table, directly ------------------------------- */ + { + int live1[1] = {0}, want1[1] = {0}, want2[2] = {0, 1}, wantB[1] = {1}; + ck(coli_cuda_init_disposition(0, 1, want1, live1) == COLI_CUDA_INIT_BUILD, + "nothing live -> BUILD"); + ck(coli_cuda_init_disposition(1, 1, want1, live1) == COLI_CUDA_INIT_ACCEPT, + "same set -> ACCEPT"); + ck(coli_cuda_init_disposition(1, 2, want2, live1) == COLI_CUDA_INIT_REFUSE, + "widening set -> REFUSE"); + ck(coli_cuda_init_disposition(1, 1, wantB, live1) == COLI_CUDA_INIT_REFUSE, + "same size, different device -> REFUSE"); + int live2[2] = {0, 1}; + ck(coli_cuda_init_disposition(2, 1, want1, live2) == COLI_CUDA_INIT_REFUSE, + "narrowing set -> REFUSE"); + ck(coli_cuda_init_disposition(2, 2, want2, live2) == COLI_CUDA_INIT_ACCEPT, + "same two-device set -> ACCEPT"); + } + + /* --- the upload gate ------------------------------------------------ */ + for (int fmt = -3; fmt <= 9; fmt++) { + if (fmt == 8) continue; + ck(coli_cuda_fp8_gate_admits(fmt, 0) && coli_cuda_fp8_gate_admits(fmt, 1), + "non-fmt=8 is admitted regardless of the LUT flag"); + } + ck(!coli_cuda_fp8_gate_admits(8, 0), "fmt=8 refused while the LUT is unpublished"); + ck(coli_cuda_fp8_gate_admits(8, 1), "fmt=8 admitted once the LUT is published"); + + /* --- the three lifecycle edges, as sequences ------------------------ */ + { + Gate g = {0, 0, {0}}; + int d0[1] = {0}, d01[2] = {0, 1}; + + ck(g_init(&g, d0, 1) == 1, "first init succeeds"); + ck(!g_upload(&g, 8), "fmt=8 refused before the first publish"); + ck(g_upload(&g, 1), "fmt=1 unaffected by the gate"); + g_set_lut(&g); + ck(g_upload(&g, 8), "fmt=8 admitted after publish"); + + /* Edge 2: same-set re-init keeps the flag, because it keeps the + * contexts the table was published to. */ + ck(g_init(&g, d0, 1) == 1, "same-set re-init returns success"); + ck(g.nctx == 1, "same-set re-init leaves the context count alone"); + ck(g_upload(&g, 8), "fmt=8 still admitted after a same-set re-init"); + + /* Edge 3: different-set re-init is refused and changes nothing. */ + ck(g_init(&g, d01, 2) == 0, "widening re-init is refused"); + ck(g.nctx == 1, "refused re-init leaves the context set alone"); + ck(g_upload(&g, 8), "refused re-init leaves the LUT flag alone"); + + /* Edge 1: shutdown clears both, and the next span must republish. */ + g_shutdown(&g); + ck(g.nctx == 0 && g.ready == 0, "shutdown clears the set and the flag"); + ck(g_init(&g, d01, 2) == 1, "after shutdown a WIDER set may be built"); + ck(!g_upload(&g, 8), "fmt=8 refused on the widened set until it republishes"); + g_set_lut(&g); + ck(g_upload(&g, 8), "fmt=8 admitted on the widened set after republish"); + } + + /* --- the invariant, over every reachable transition sequence -------- */ + { + /* Exhaustive over sequences of length 6 drawn from + * {init{0}, init{0,1}, set_lut, shutdown}: no reachable state may have + * a live-free context set and a set flag, and no state may admit fmt=8 + * without a live set. */ + int d0[1] = {0}, d01[2] = {0, 1}; + int idx[6] = {0, 0, 0, 0, 0, 0}; + long checked = 0; + for (;;) { + Gate g = {0, 0, {0}}; + for (int s = 0; s < 6; s++) { + switch (idx[s]) { + case 0: g_init(&g, d0, 1); break; + case 1: g_init(&g, d01, 2); break; + case 2: g_set_lut(&g); break; + default: g_shutdown(&g); break; + } + if (g.nctx == 0 && g.ready != 0) { printf("FAIL invariant: no contexts but the LUT flag is set\n"); return 1; } + if (g.nctx == 0 && g_upload(&g, 8)) { printf("FAIL invariant: fmt=8 admitted with no live context\n"); return 1; } + } + checked++; + int p = 5; + while (p >= 0 && ++idx[p] > 3) { idx[p] = 0; p--; } + if (p < 0) break; + } + if (checked != 4096) { printf("FAIL sequence enumeration covered %ld, expected 4096\n", checked); return 1; } + printf("lut-gate invariant holds over %ld transition sequences\n", checked); + } + + if (fails) { printf("cuda lut-gate tests: %d FAILED\n", fails); return 1; } + printf("cuda lut-gate tests: ok\n"); + return 0; +} diff --git a/c/tests/test_fp8_cuda.cu b/c/tests/test_fp8_cuda.cu index 87ccecc26..fe374fd06 100644 --- a/c/tests/test_fp8_cuda.cu +++ b/c/tests/test_fp8_cuda.cu @@ -335,5 +335,55 @@ int main(void){ } coli_cuda_shutdown(); } + + /* ---- Phase 3: LUT-gate lifecycle across shutdown/re-init --------------- */ + { /* g_fp8_lut_ready is process-wide while the e4m3 table is per-device. + * SHUTDOWN is the only site that clears it; coli_cuda_init never writes + * it. That is sufficient because init refuses to rebuild contexts while + * a set is live: a re-init naming the SAME set returns 1 and leaves the + * contexts -- and therefore the published table -- untouched, and one + * naming a DIFFERENT set is refused outright. So the device set cannot + * widen past what the last publish covered without passing through + * shutdown, which clears the flag. This block pins all three edges. */ + int devs[1]={0}; + enum { LO=4, LI=128 }; + uint8_t lw[LO*LI]; float ls[1]={1.f}; + for(size_t i=0;i1 ? 1 : ndev}; /* {0,1} on a multi-GPU box, else out of range */ + if(coli_cuda_init(other,2)!=0){ printf("FAIL different-set re-init was not refused\n"); return 1; } + if(coli_cuda_device_count()!=1){ printf("FAIL refused re-init still changed the context set\n"); return 1; } + if(!coli_cuda_tensor_upload(<,lw,ls,8,LI,LO,0)){ printf("FAIL refused re-init disturbed the LUT flag\n"); return 1; } + coli_cuda_tensor_free(lt); lt=nullptr; + } + + coli_cuda_shutdown(); + printf("lut-gate lifecycle: shutdown clears; same-set re-init keeps; different-set re-init refused\n"); + } printf("OK\n"); return 0; } diff --git a/c/tests/test_shard_kvb_refuse.c b/c/tests/test_shard_kvb_refuse.c new file mode 100644 index 000000000..ff6cadc42 --- /dev/null +++ b/c/tests/test_shard_kvb_refuse.c @@ -0,0 +1,151 @@ +/* layer_cuda_shard_kvb() (colibri.c, COLI_CUDA_ATTN_SHARD) -- the multi-device kv_b + * head-shard uploader. Its rb/weights/scale arithmetic is written for exactly fmt=1/2/3/4 + * (per-row byte strides, per-row or per-group scales); before the guard this file proves, + * an fmt=8 (fp8-e4m3-b128) kv_b was ADMITTED and only failed safe BY ACCIDENT: fmt=8 + * keeps its raw e4m3 bytes in q8 (q4 stays NULL, per the QT struct comment), so the + * function selected a NULL weight pointer with an int2 row stride and + * coli_cuda_tensor_upload_g's !weights check happened to reject the upload before + * anything dereferenced it -- silent, unnamed, and one refactor away from a misread + * (fmt=8's per-128x128-BLOCK scale array would ALSO have been sliced with per-row + * geometry). This probe pins the explicit refusal that replaced the accident: + * (1) an fmt=8 kv_b shard attempt refuses BY NAME on stderr, BEFORE any pointer/stride + * use, and the message says what serves fmt=8 instead (the absorb path on the + * layer home device); + * (2) no shard state is minted (n_kv_b_shard==0, kv_b_shard[] untouched, + * kv_b.cuda_eligible unchanged); + * (3) the notice is bounded (once per process per fmt, never per layer); + * (4) other un-shardable fmts (fmt=6 here) refuse by name too, with their own message; + * (5) an allowlisted fmt (fmt=1) is NOT refused -- with no device context initialized + * its upload fails silently and no shard is minted, but no refusal line appears, + * so the guard is format-targeted, not a blanket gate. + * PROOF-OF-BITE: built against the pre-guard colibri.c, (1)/(3)/(4) fail -- the fmt=8 + * call slid past the format check into the accidental-safe upload rejection with no + * message at all. Needs -DCOLI_CUDA (CUDA=1) to compile the function; a CPU-only build + * SKIPs loudly instead of pretending to cover it. No GPU work is performed: every path + * exercised here returns before any device context exists, so this runs on a CUDA build + * host even without a card. + * Portable stderr capture: freopen/dup2, same seam as tests/test_kvb_notice.c. */ +#define main coli_glm_main_unused +#include "../colibri.c" +#undef main + +#include +#include +#include + +#ifndef COLI_CUDA +int main(void){ + printf("test_shard_kvb_refuse: SKIP (built without COLI_CUDA -- build with CUDA=1 on a CUDA host to exercise layer_cuda_shard_kvb)\n"); + return 0; +} +#else + +#include + +static int fails = 0; +#define CHECK(c) do{ if(!(c)){ printf("FAIL %s:%d: %s\n", __FILE__, __LINE__, #c); fails++; } }while(0) + +static int redirect_stderr(const char *path){ + fflush(stderr); + int saved = dup(fileno(stderr)); + if(saved<0){ printf("FAIL: dup(stderr) failed -- capture seam unusable, aborting\n"); exit(1); } + if(!freopen(path, "w+", stderr)){ + printf("FAIL: freopen(%s) failed -- capture seam unusable, aborting\n", path); + exit(1); + } + return saved; +} +static void restore_stderr(int saved, char *buf, size_t bufsz){ + fflush(stderr); + long n = ftell(stderr); + if(n<0) n=0; + rewind(stderr); + size_t want = (size_t)n < bufsz-1 ? (size_t)n : bufsz-1; + size_t got = fread(buf, 1, want, stderr); + buf[got] = 0; + fflush(stderr); + dup2(saved, fileno(stderr)); + close(saved); +} + +static int count_sub(const char *hay, const char *needle){ + int n=0; const char *p=hay; + while((p=strstr(p,needle))){ n++; p++; } + return n; +} + +/* kv_b-shaped fixture at H=4, Q=8, V=8 -> O=H*(Q+V)=64; I=200 (nblkI=2, partial column + * tail -- the exact scale geometry the shard's per-row slicing would have misread). */ +enum { H=4, Q=8, V=8, O=H*(Q+V), I=200 }; + +static void mk_layer_fmt8(Layer *l, int8_t *q8, float *s){ + memset(l,0,sizeof *l); + for(int64_t i=0;i<(int64_t)O*I;i++) q8[i]=(int8_t)(i*37+11); + int64_t nblk=fp8_nblk(O)*fp8_nblk(I); + for(int64_t i=0;ikv_b.fmt=8; l->kv_b.O=O; l->kv_b.I=I; l->kv_b.gs=0; + l->kv_b.q8=q8; l->kv_b.s=s; /* q4 stays NULL: the fmt=8 convention */ + l->kv_b.cuda_eligible=1; l->kv_b.cuda_device=0; +} + +static void run_shard(Layer *l, char *buf, size_t bufsz){ + int saved = redirect_stderr("tests/tmp_shard_kvb_refuse.stderr"); + layer_cuda_shard_kvb(l,H,Q,V); + restore_stderr(saved, buf, bufsz); + remove("tests/tmp_shard_kvb_refuse.stderr"); +} + +int main(void){ + /* pretend a 2-device dense-CUDA setup so the function's early gate passes; no + * device context is ever initialized, and none is needed (see file header). */ + g_cuda_enabled=1; g_cuda_dense=1; g_cuda_ndev=2; + g_cuda_devices[0]=0; g_cuda_devices[1]=0; + + static int8_t q8[(int64_t)O*I]; + static float s[((O+127)/128)*((I+127)/128)]; + static float srow[O]; /* per-row scales for the allowlisted fmt=1 case */ + char err[4096]; + Layer l; + + /* (1)+(2): fmt=8 refuses by name, names the absorb path, mints no shard state */ + mk_layer_fmt8(&l,q8,s); + run_shard(&l,err,sizeof err); + CHECK(strstr(err,"layer_cuda_shard_kvb")!=NULL); + CHECK(strstr(err,"refus")!=NULL); + CHECK(strstr(err,"fmt=8")!=NULL); + CHECK(strstr(err,"absorb path")!=NULL); /* says what serves fmt=8 instead */ + CHECK(l.n_kv_b_shard==0); + CHECK(l.kv_b_shard[0]==NULL && l.kv_b_shard[1]==NULL); + CHECK(l.kv_b.cuda_eligible==1); /* shard bookkeeping never ran */ + + /* (3): bounded -- a second fmt=8 layer (any of the other 60) adds no second line */ + mk_layer_fmt8(&l,q8,s); + run_shard(&l,err,sizeof err); + CHECK(count_sub(err,"layer_cuda_shard_kvb")==0); + CHECK(l.n_kv_b_shard==0); + + /* (4): fmt=6 (E8/IQ3, single 4-byte scale tag) refuses by name with its own message */ + memset(&l,0,sizeof l); + l.kv_b.fmt=6; l.kv_b.O=O; l.kv_b.I=I; l.kv_b.gs=0; + l.kv_b.q4=(uint8_t*)q8; l.kv_b.s=s; /* non-NULL on purpose: only the guard saves it */ + run_shard(&l,err,sizeof err); + CHECK(strstr(err,"layer_cuda_shard_kvb")!=NULL); + CHECK(strstr(err,"refus")!=NULL); + CHECK(strstr(err,"fmt=6")!=NULL); + CHECK(l.n_kv_b_shard==0); + + /* (5): allowlisted fmt=1 is NOT refused -- upload fails silently (no device + * context), no shard minted, but no refusal line either */ + memset(&l,0,sizeof l); + for(int i=0;ikv_b) carries explicit fmt==8 decode arms, so a container minted here -loads clean (qt_from_disk resolves the stamp exactly like o_proj's) and -decodes correctly through the batched (`kvs`-nonNULL) serving path, where -absorb cannot be bypassed. The CUDA absorb kernels still have no fmt==8 -case (absorb_fmt_ok delegates to coli_cuda_weight_at_supported's fmt 0..4 -allowlist), so fmt=8 kv_b decode stays on the CPU absorb path until that -lands. +l->kv_b) carries explicit fmt==8 decode arms, and the CUDA absorb kernels +decode fmt=8 through weight_at/absorb_scale (admitted by +coli_cuda_weight_at_supported, uploads gated on the published e4m3 LUT). +A container minted here loads clean (qt_from_disk resolves the stamp +exactly like o_proj's) and decodes correctly through the batched +(`kvs`-nonNULL) serving path, where absorb cannot be bypassed. METADATA STAMP (reference implementation of the FORMATS-registry FR -- see docs/FORMATS.md): every output shard's safetensors `__metadata__` carries a From ea8a7e4d55d9c538d9ef7d6a8fadfbe94c84ebff Mon Sep 17 00:00:00 2001 From: monotophic Date: Thu, 3 Sep 2026 15:06:00 -0400 Subject: [PATCH 4/6] test(fp8): e2e serve-batch pin, refusal canary, and mint-to-load regression test_fp8_serve_batch_e2e.py pins the fmt=8 kv_b batched-serve decode end to end against a real container, and is a named SKIP when none is present. Its child environment is allowlist-scrubbed: ambient engine knobs are stripped, read-only model-location settings pass, documented legacy aliases are stripped alongside their primaries, and the ratified non-absorb arm passes value-restricted. The CUDA absorb path is opt-in: the test INJECTS the exact engine bundle that path needs rather than letting it arrive ambiently, then witnesses the engine's boot report and fails outright on the resident-dense-on-CPU line, so a misconfigured lane cannot bank a vacuous green. That witness caught a first attempt which value-allowed the bundle through the scrub instead of injecting it. test_fp8_refusal_canary.py runs in CI without a container: it pins that each absorb function keeps at least one refusal matching the end-to-end matcher -- existential on purpose, since an additional differently-worded guard is new coverage rather than drift -- and that the matcher stays selective against both a synthetic near-miss and the real in-tree sibling. test_e8x4g64_mint_load.py and test_e8x4g64_loader.c pin the container class through the real conversion CLI and the real C loader at toy scale, plus the duplicate-tensor-name refusal on a tool-produced container. The C harness takes a container directory on argv, so it is excluded from the gate list and driven by the Python test; it gains a Makefile rule anyway, so that it compiles under the suite's own flags instead of only the driver's hand-copied list, which is what its no-warnings assertion had been vouching for. The driver no longer swallows a failed OpenMP probe on macOS: silently compiling a different binary from the one the Makefile builds is what made that assertion misleading. The external-files table in the repack tool now cites symbols instead of line numbers. Those anchors had rotted twice while the table claimed to be verified against the current tree, and a stale line number reads as a verification that was not performed. --- c/Makefile | 15 +- c/tests/test_e8x4g64_loader.c | 108 +++++ c/tests/test_e8x4g64_mint_load.py | 297 +++++++++++++ c/tests/test_fp8_refusal_canary.py | 120 +++++ c/tests/test_fp8_serve_batch_e2e.py | 668 ++++++++++++++++++++++++++++ c/tools/repack_fp8_passthrough.py | 20 +- docs/FORMATS.md | 5 +- 7 files changed, 1221 insertions(+), 12 deletions(-) create mode 100644 c/tests/test_e8x4g64_loader.c create mode 100644 c/tests/test_e8x4g64_mint_load.py create mode 100644 c/tests/test_fp8_refusal_canary.py create mode 100644 c/tests/test_fp8_serve_batch_e2e.py diff --git a/c/Makefile b/c/Makefile index 1ee880c11..eed827041 100644 --- a/c/Makefile +++ b/c/Makefile @@ -498,12 +498,17 @@ TEST_RULES := $(shell sed -n 's|^tests/\(test_[a-z0-9_]*\)\$$(EXE):.*|\1|p' $(f # "unsupported option '-msse4.1' for target arm64-...". # Appended back below only on x86-64 hosts, alongside the other conditional # platform tests already handled the same way. +# test_e8x4g64_loader takes a minted container directory on argv and prints its +# usage (exit 2) without one, so it is a harness rather than a gate: it has a +# rule here so it compiles under the suite's own $(CFLAGS) and its warnings are +# caught, and tests/test_e8x4g64_mint_load.py is what actually drives it. # test_qwen38_tier_engine drives qwen38.c over the generated FP8 fixture # (qwen38_tiny_fp8, gitignored); it runs from qwen38-tier-engine-check. TEST_EXCLUDE = test_uring test_deepseek_v4 test_v4_ownership test_v4_serve_framing \ test_segment_adapters_registration test_segment_adapters_real \ test_edge_adapters_registration test_edge_adapters_real \ - test_gsgemv_sse41 test_qgemv_sse41 test_olmoe_dot_i8_16_sse41 test_qwen38_tier_engine + test_gsgemv_sse41 test_qgemv_sse41 test_olmoe_dot_i8_16_sse41 test_qwen38_tier_engine \ + test_e8x4g64_loader TEST_BINS = $(addprefix tests/,$(addsuffix $(EXE),$(filter-out $(TEST_EXCLUDE),$(TEST_RULES)))) ifneq (,$(LINUX)) TEST_BINS += tests/test_uring$(EXE) @@ -2096,6 +2101,14 @@ tests/test_olmoe_dot_i8_16_sse41$(EXE): tests/test_olmoe_dot_i8_16.c olmoe.c sse tests/test_shard_kvb_refuse$(EXE): tests/test_shard_kvb_refuse.c colibri.c st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h quant.h fp8_format.h sample.h kv_persist.h telemetry.h $(CUDA_OBJ) $(CC) $(CFLAGS) $< $(CUDA_OBJ) -o $@ $(LDFLAGS) +# The e8x4g64 container loader harness. It takes a minted directory on argv and +# is driven by tests/test_e8x4g64_mint_load.py, which builds it too -- this rule +# exists so it also compiles under the suite's own $(CFLAGS) rather than only +# under the driver's hand-copied flag list, and so a warning regression in it +# fails the normal build. +tests/test_e8x4g64_loader$(EXE): tests/test_e8x4g64_loader.c st.h quant.h fp8_format.h compat.h + $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) + test-c: $(TEST_BINS) $(PYTHON) tools/run_tests.py $(TEST_BINS) diff --git a/c/tests/test_e8x4g64_loader.c b/c/tests/test_e8x4g64_loader.c new file mode 100644 index 000000000..32fb17072 --- /dev/null +++ b/c/tests/test_e8x4g64_loader.c @@ -0,0 +1,108 @@ +/* wq-v0-class (e8x4g64: int8 spine + grouped-int4 g64 experts) END-TO-END + * loader harness -- the C-side half of the mint->load regression case. + * + * tests/test_e8x4g64_mint_load.py drives it: builds a synthetic FP8 GLM-shaped + * checkpoint (tools/glm_fp8_emit.py, the same fixture helper the fp8 + * passthrough e2e uses), mints it with the REAL tools/convert_fp8_to_int4.py + * CLI at the e8x4g64 recipe (--ebits 8 --xbits 4 --io-bits 8 --group-size 64), + * then invokes this binary against the real output directory. Sibling of + * tests/test_fp8_e2e_loader.c (fmt=8 passthrough), same division of labor: + * the Python side proves the TOOL ran for real; this side proves the REAL + * st_init/qt_from_disk (the identical functions every model load uses) + * resolve every minted tensor to the format the recipe promises. + * + * Usage: test_e8x4g64_loader [name O I wantfmt wantgs]... + * For every (name,O,I,wantfmt,wantgs) 5-tuple, calls the REAL qt_from_disk and + * asserts: + * (a) the resolved fmt equals wantfmt (1 = int8 per-row spine, 4 = grouped + * int4 experts) and, for fmt=4, the derived group size equals wantgs -- + * both resolved by qt_resolve_fmt's byte arithmetic on the tool's actual + * output, neither mocked nor stamped; + * (b) the weight and scale buffers are non-NULL; + * (c) every dequantized value is finite -- catches a scale-layout or + * packing bug a pure byte-count check wouldn't. + * + * The D-2 duplicate-name refusal is exercised through this same binary: the + * Python driver runs it a second time against a copy of the container holding + * a duplicated shard, and st_init below then refuses (exit 1, naming both + * shards) before any tuple is checked -- the positive half of the D-2 pair, + * on a tool-produced container. The clean run doubles as the negative + * control: a legitimate mint loads with no refusal. */ +#define main coli_glm_main_unused +#include "../colibri.c" +#undef main + +#include +#include +#include +#include + +int main(int argc, char **argv){ + if(argc < 2){ fprintf(stderr,"usage: %s [name O I wantfmt wantgs]...\n", argv[0]); return 2; } + if((argc-2) % 5 != 0){ + fprintf(stderr,"args after must come in (name,O,I,wantfmt,wantgs) 5-tuples (got %d)\n", argc-2); + return 2; + } + const char *dir = argv[1]; + int ntensors = (argc-2)/5; + if(ntensors == 0){ fprintf(stderr,"no tuples given -- nothing to check\n"); return 2; } + + static Model gm; memset(&gm,0,sizeof gm); + st_init(&gm.S, dir); /* D-2 duplicate-name detection lives in here (st.h) */ + + int fails = 0; + for(int i=0;i>1]; int nib=(ii&1)?(b>>4):(b&0xF); + float v=((int)nib-8)*scl[ii/t.gs]; + if(!isfinite(v)){ bad=o*(int64_t)I+ii; break; } } } + } else { + printf("FAIL %s: this harness only knows the e8x4g64 formats (1, 4); " + "wantfmt=%d is a driver bug\n", name, wantfmt); + fails++; continue; + } + if(bad >= 0){ + printf("FAIL %s: non-finite dequantized value at flat index %lld\n", name, (long long)bad); + fails++; continue; + } + printf("ok %s: fmt=%d%s O=%d I=%d, loaded through the real loader, all-finite\n", + name, t.fmt, t.fmt==4?" (g64)":"", O, I); + } + if(fails){ printf("e8x4g64 mint->load: %d/%d tensor(s) FAILED\n", fails, ntensors); return 1; } + printf("e8x4g64 mint->load: ok (%d tensor(s))\n", ntensors); + return 0; +} diff --git a/c/tests/test_e8x4g64_mint_load.py b/c/tests/test_e8x4g64_mint_load.py new file mode 100644 index 000000000..7f506950a --- /dev/null +++ b/c/tests/test_e8x4g64_mint_load.py @@ -0,0 +1,297 @@ +"""wq-v0-class regression case: REAL tools/convert_fp8_to_int4.py output fed +into the REAL C loader (st_init/qt_from_disk in colibri.c) at toy scale. + +The wq-v0 production container class is "e8x4g64": int8 per-row spine +(attention / shared expert / dense MLP / embed / lm_head -> fmt=1) with +grouped-int4 g64 routed experts (fmt=4, gs=64). This test mints that class +each cycle at toy scale with GLM-shaped tensor names and the load-bearing +real dimensions (kv_lora_rank-width kv_b contraction, group-64-divisible +expert dims) through the real converter CLI -- not a reimplementation of its +quantizers -- and then proves the REAL C loader resolves every tensor to the +format the recipe promises. Regression coverage, not new-feature proof: the +full-scale wq-v0 container is already banked and audited; what this pins is +that the TOOLING and the LOADER still agree on the class after engine +changes (the mint half of the mint->load->run regression bar; the run half +executes at full scale against the banked container, outside this suite). + +Also carried here, because it belongs to exactly this container class: the +D-2 duplicate-tensor-name pair on a TOOL-PRODUCED container. + * positive: a copy of the minted container with one shard duplicated must + refuse at st_init with the exact D-2 message (naming both shards), before + any tensor is read; + * negative control: the untouched mint loads clean through the same + binary -- the guard is duplicate-targeted, not a blanket gate. +(tests/test_dup_name_refusal.c pins the same guard on hand-built fixtures; +this is the tool-produced end of it.) + +Hermetic: checkpoint, minted output, duplicated copy, and the compiled +harness all live under temporary directories; nothing is left behind. +""" +import glob +import os +import shutil +import subprocess +import sys +import tempfile +import unittest + +try: + import torch +except ImportError as e: + raise unittest.SkipTest(f"torch not installed: {e}") + +try: + import numpy as np +except ImportError as e: + raise unittest.SkipTest(f"numpy not installed: {e}") + +HERE = os.path.dirname(os.path.abspath(__file__)) +C_DIR = os.path.normpath(os.path.join(HERE, "..")) +sys.path.insert(0, os.path.join(C_DIR, "tools")) +# reuse: the real fp8 block-quantize fixture helpers, not reimplemented. The +# quantize/dequantize pair anchors the numeric reference below to EXACTLY what +# the converter ingests (its dequant() applies the same repeat_interleave +# formula to the same saved fp8 bytes); the converter's own int8/int4 +# quantizers are NOT imported anywhere in this file -- their output is decoded +# and bounded by independent code below. +from glm_fp8_emit import save_fp8_safetensors, fp8_block_quantize, fp8_block_dequantize + + +def _cc_flags(): + """Mirror the Makefile's CFLAGS closely enough to compile colibri.c + cleanly (same arrangement as tests/test_fp8_e2e_repack_load.py).""" + cc = shutil.which("cc") or shutil.which("clang") or shutil.which("gcc") + if not cc: + return None, None, None + cflags = ["-O3", "-Wall", "-Wextra", "-Wno-unused-parameter", + "-Wno-misleading-indentation", "-Wno-unused-function"] + ldflags = ["-lm"] + if sys.platform not in ("darwin", "win32"): + cflags += ["-fopenmp"] + ldflags += ["-fopenmp"] + if sys.platform == "darwin": + # The Makefile builds this harness with OpenMP on macOS via Homebrew's + # libomp. Locating it must not fail quietly: a swallowed error here + # compiles a DIFFERENT binary from the one the suite builds, and the + # warning assertion below would then vouch for flags nobody ships. + # Report why it could not be found and let the caller decide. + try: + probe = subprocess.run(["brew", "--prefix", "libomp"], capture_output=True, + text=True, timeout=10) + prefix = probe.stdout.strip() + if probe.returncode != 0: + raise RuntimeError( + "`brew --prefix libomp` exited %d: %s" + % (probe.returncode, probe.stderr.strip() or "(no stderr)")) + except FileNotFoundError: + raise RuntimeError( + "libomp is required to build this harness on macOS the way the " + "Makefile builds it, and `brew` is not on PATH") + except (OSError, subprocess.TimeoutExpired) as exc: + raise RuntimeError("could not run `brew --prefix libomp`: %r" % (exc,)) + inc, lib = os.path.join(prefix, "include"), os.path.join(prefix, "lib") + if not (prefix and os.path.exists(os.path.join(inc, "omp.h"))): + raise RuntimeError( + "`brew --prefix libomp` gave %r but %s/omp.h does not exist; " + "install libomp (`brew install libomp`) rather than compiling " + "this harness without OpenMP" % (prefix, inc)) + cflags += ["-Xclang", "-fopenmp", "-I", inc] + ldflags += ["-L", lib, "-lomp"] + return cc, cflags, ldflags + + +# Toy scale, real geometry where it is load-bearing: +# * kv_b_proj keeps the REAL contraction width I=512 (GLM-5.2's +# kv_lora_rank) with a toy head count -> O=160; +# * routed-expert dims stay multiples of 64 (real containers are), so the +# g64 grouping has no synthetic tail the production mint never sees -- +# while D_H (o_proj/q_a rows, expert contraction) is NOT a multiple of +# 128, keeping the fp8 input's block scales on the partial-tail path; +# * no tensor sits on the fmt=1-vs-fmt=8 collision boundary +# (O == ceil(O/128)*ceil(I/128)), so byte-arithmetic resolution is +# unambiguous, as it is at full scale. +D_H = 320 # toy hidden size (multiple of 64, not of 128) +KV_B_O, KV_B_I = 160, 512 +E_M = 192 # toy expert intermediate (multiple of 64) +VOCAB = 512 + + +class E8x4g64MintLoadTest(unittest.TestCase): + """The real e8x4g64 mint recipe -> the real C loader.""" + + _harness_dir = None + _harness_bin = None + _build_stderr = None + + @classmethod + def setUpClass(cls): + cc, cflags, ldflags = _cc_flags() + if not cc: + raise unittest.SkipTest("no C compiler found on PATH") + cls._harness_dir = tempfile.TemporaryDirectory() + harness_src = os.path.join(HERE, "test_e8x4g64_loader.c") + harness_bin = os.path.join(cls._harness_dir.name, "test_e8x4g64_loader") + build = subprocess.run([cc] + cflags + [harness_src, "-o", harness_bin] + ldflags, + capture_output=True, text=True, cwd=C_DIR) + if build.returncode != 0: + raise AssertionError(f"e8x4g64 harness build failed:\n{build.stderr}") + cls._harness_bin = harness_bin + cls._build_stderr = build.stderr + + @classmethod + def tearDownClass(cls): + if cls._harness_dir: + cls._harness_dir.cleanup() + + def setUp(self): + self.tmp = tempfile.TemporaryDirectory() + self.indir = os.path.join(self.tmp.name, "fp8src") + os.makedirs(self.indir) + self.outdir = os.path.join(self.tmp.name, "out") + self.shard = os.path.join(self.indir, "model-00001-of-00001.safetensors") + + def tearDown(self): + self.tmp.cleanup() + + def _emit_checkpoint(self): + """A GLM-shaped fp8 source checkpoint: every name below is one the real + converter's classify() routes exactly like the production checkpoint's + (kvb/attn/o -> ebits, sh/dmlp -> ebits, io -> io_bits, experts -> xbits, + norms -> f32 passthrough).""" + torch.manual_seed(11) + L = "model.layers.0" + sd = { + f"{L}.self_attn.kv_b_proj.weight": torch.randn(KV_B_O, KV_B_I) * 0.02, + f"{L}.self_attn.q_a_proj.weight": torch.randn(D_H, D_H) * 0.02, + f"{L}.self_attn.o_proj.weight": torch.randn(D_H, D_H) * 0.02, + f"{L}.mlp.shared_experts.gate_proj.weight": torch.randn(E_M, D_H) * 0.02, + f"{L}.mlp.experts.0.gate_proj.weight": torch.randn(E_M, D_H) * 0.02, + f"{L}.mlp.experts.0.up_proj.weight": torch.randn(E_M, D_H) * 0.02, + f"{L}.mlp.experts.0.down_proj.weight": torch.randn(D_H, E_M) * 0.02, + "model.embed_tokens.weight": torch.randn(VOCAB, D_H) * 0.02, + f"{L}.input_layernorm.weight": torch.randn(D_H), # f32 passthrough control + } + self.sd = sd # kept: the numeric round-trip test derives its reference from it + save_fp8_safetensors(sd, self.shard) + + def _mint(self): + """The REAL converter CLI at the wq-v0 e8x4g64 recipe: int8 spine + (--ebits 8, --io-bits 8), grouped-int4 g64 routed experts (--xbits 4, + --group-size 64).""" + tool = os.path.join(C_DIR, "tools", "convert_fp8_to_int4.py") + rc = subprocess.run([sys.executable, tool, "--indir", self.indir, + "--outdir", self.outdir, "--n-layers", "1", + "--ebits", "8", "--xbits", "4", "--io-bits", "8", + "--group-size", "64"], + capture_output=True, text=True) + self.assertEqual(rc.returncode, 0, + f"real converter failed:\nSTDOUT:\n{rc.stdout}\nSTDERR:\n{rc.stderr}") + outs = glob.glob(os.path.join(self.outdir, "out-*.safetensors")) + self.assertEqual(len(outs), 1, f"expected exactly one minted shard, got {outs}") + return outs[0] + + # 5-tuples for the C harness: (name, O, I, wantfmt, wantgs) + _EXPECT = [ + ("model.layers.0.self_attn.kv_b_proj.weight", KV_B_O, KV_B_I, 1, 0), + ("model.layers.0.self_attn.q_a_proj.weight", D_H, D_H, 1, 0), + ("model.layers.0.self_attn.o_proj.weight", D_H, D_H, 1, 0), + ("model.layers.0.mlp.shared_experts.gate_proj.weight", E_M, D_H, 1, 0), + ("model.embed_tokens.weight", VOCAB, D_H, 1, 0), + ("model.layers.0.mlp.experts.0.gate_proj.weight", E_M, D_H, 4, 64), + ("model.layers.0.mlp.experts.0.up_proj.weight", E_M, D_H, 4, 64), + ("model.layers.0.mlp.experts.0.down_proj.weight", D_H, E_M, 4, 64), + ] + + def _load(self, container_dir): + args = [self._harness_bin, container_dir] + for name, n_out, n_in, fmt, gs in self._EXPECT: + args += [name, str(n_out), str(n_in), str(fmt), str(gs)] + return subprocess.run(args, capture_output=True, text=True) + + def test_harness_builds_without_warnings(self): + """Production flags require zero warnings (pr-cycle mechanical gate).""" + self.assertEqual(self._build_stderr.strip(), "", + f"e8x4g64 harness build produced warnings:\n{self._build_stderr}") + + def test_e8x4g64_mint_loads_through_real_c_loader(self): + self._emit_checkpoint() + self._mint() + rc = self._load(self.outdir) + self.assertEqual(rc.returncode, 0, + f"loader harness failed:\nSTDOUT:\n{rc.stdout}\nSTDERR:\n{rc.stderr}") + for name, _, _, fmt, _ in self._EXPECT: + self.assertIn(f"ok {name}: fmt={fmt}", rc.stdout) + # negative D-2 control rides along: the clean mint produced NO + # duplicate-name refusal anywhere in the load + self.assertNotIn("duplicate tensor name", rc.stderr) + + def test_e8x4g64_numeric_round_trip(self): + """The minted VALUES, not just the format class (deep-audit MAJOR-1: + the fmt/finite checks alone passed a 37x scale corruption of + quant_int8 -- format class and finiteness both survive numeric + corruption). This closes DR-9(a)'s literal words, "mint ROUND-TRIP at + toy scale": dequantize the minted bytes with independent numpy code + (the converter's quantizers are never imported here) and require + agreement with the fp8-round-tripped source within the schemes' own + exact bound -- symmetric absmax rint quantization can never miss by + more than half a quantization step, so the tolerance is 0.5 steps + (+1e-3 for f32 arithmetic slop), derived from a scale RECOMPUTED from + the source, never from the minted .qs (a corrupted stored scale must + widen the error, not the bound).""" + self._emit_checkpoint() + minted = self._mint() + from safetensors.numpy import load_file + out = load_file(minted) + for name, n_out, n_in, fmt, gs in self._EXPECT: + # Reference: exactly what the converter ingested -- the fp8 + # block-quantized source, dequantized (bit-identical formula to + # the converter's dequant(); see the import comment above). + w_fp8, scale = fp8_block_quantize(self.sd[name].float()) + w_ref = fp8_block_dequantize(w_fp8, scale).numpy().astype(np.float32) + self.assertEqual(w_ref.shape, (n_out, n_in)) + q, qs = out[name], out[name + ".qs"] + if fmt == 1: + # int8 per-row: one scale per output row, q in [-127,127] + # (amax maps to qmax=127, so -128 is unreachable). + deq = q.view(np.int8).reshape(n_out, n_in).astype(np.float32) * qs.reshape(n_out, 1) + s_true = np.maximum(np.abs(w_ref).max(axis=1, keepdims=True) / 127.0, 1e-8) + step = np.broadcast_to(s_true, (n_out, n_in)) + else: + # int4 grouped g64: packed nibbles (low = even column), one + # scale per (row, 64-column group), q in [-7,7] likewise. + rb, ng = (n_in + 1) // 2, (n_in + gs - 1) // gs + b = q.reshape(n_out, rb) + deq = np.empty((n_out, n_in), np.float32) + deq[:, 0::2] = (b & 0xF).astype(np.float32)[:, : (n_in + 1) // 2] - 8.0 + deq[:, 1::2] = (b >> 4).astype(np.float32)[:, : n_in // 2] - 8.0 + deq *= np.repeat(qs.reshape(n_out, ng), gs, axis=1)[:, :n_in] + pad = np.zeros((n_out, ng * gs - n_in), np.float32) + grp = np.concatenate([np.abs(w_ref), pad], axis=1).reshape(n_out, ng, gs) + s_true = np.maximum(grp.max(axis=2) / 7.0, 1e-8) # [n_out, ng] + step = np.repeat(s_true, gs, axis=1)[:, :n_in] + worst = float((np.abs(deq - w_ref) / step).max()) + self.assertLessEqual( + worst, 0.5 + 1e-3, + f"{name}: minted values diverge from the source by {worst:.4g} " + f"quantization steps (bound: half a step) -- the converter's " + f"numeric output regressed even though the format class may " + f"still be intact") + + def test_d2_duplicate_shard_refuses_by_name(self): + """D-2 positive on the tool-produced container: duplicating the minted + shard under a second indexed name must refuse at st_init with the + exact duplicate-tensor-name message, before any tensor check runs.""" + self._emit_checkpoint() + minted = self._mint() + dupdir = os.path.join(self.tmp.name, "dup") + shutil.copytree(self.outdir, dupdir) + shutil.copy(minted, os.path.join(dupdir, "out-99999.safetensors")) + rc = self._load(dupdir) + self.assertNotEqual(rc.returncode, 0, "duplicated container must refuse to load") + self.assertIn("duplicate tensor name across indexed shards, refusing", rc.stderr) + self.assertNotIn("ok model.", rc.stdout, + "no tensor may load from a container st_init refused") + + +if __name__ == "__main__": + unittest.main() diff --git a/c/tests/test_fp8_refusal_canary.py b/c/tests/test_fp8_refusal_canary.py new file mode 100644 index 000000000..4d6dda306 --- /dev/null +++ b/c/tests/test_fp8_refusal_canary.py @@ -0,0 +1,120 @@ +"""Drift canary for the absorb-path fmt-refusal message shape. + +tests/test_fp8_serve_batch_e2e.py machine-matches the engine's death on the +batched fmt=8 serve path with one regex (its REFUSAL constant, the family +`(qt_addrow|qt_matvec_rows): unsupported fmt=`), in two load-bearing +places: the pass lane's engine-death scan, and the expect-bite mode the +fleet's old-binary half PASSES ONLY THROUGH. The product strings live in +c/colibri.c (qt_addrow's refusal and qt_matvec_rows'); an innocent reword +of either -- "unsupported" to "unhandled", dropping the function-name +prefix -- would not break any build or unit suite, but it would silently +degrade the bite proof's machine-match into a generic "engine died", and +the e2e test only runs where a real fmt=8 container exists (never in CI). +This canary closes that gap in CI: it fails, by name, the moment the +matcher and the source strings drift apart, in either direction. + +What is load-bearing is exactly what this file asserts, no more: + * each function still has at least one refusal-class stderr message + (its runtime firing is pinned separately by tests/test_qt_addrow.c's + fork+waitpid refusal cases, which also require the "refus" word); + * at least ONE refusal-class message from each function matches the e2e + REFUSAL regex -- imported from the e2e module, never re-typed, so the + two files cannot agree by coincidence (existential, not universal: a + future second, differently-worded guard in the same function is a new + refusal, not drift, and must not fail this canary); + * the regex is still SELECTIVE -- it must not match an arbitrary death + line, or expect-bite would file any crash as the bite. +Everything else about the messages (wording after the matched prefix, +fmt lists, line breaks) is deliberately unpinned: message edits that keep +the matchable shape must stay free. +""" +import re +import sys +import unittest +from pathlib import Path + +HERE = Path(__file__).resolve().parent +sys.path.insert(0, str(HERE)) # sibling import under any invocation style +from test_fp8_serve_batch_e2e import REFUSAL + +COLIBRI_C = HERE.parent / "colibri.c" +FUNCTIONS = ("qt_addrow", "qt_matvec_rows") + +# One C string literal (escapes included), and an fprintf-to-stderr whose +# message is one or more adjacent literals (the refusals wrap across lines). +_STRING = r'"(?:[^"\\\n]|\\.)*"' +_FPRINTF = re.compile(r'fprintf\s*\(\s*stderr\s*,\s*(' + _STRING + r'(?:\s*' + _STRING + r')*)') + + +def stderr_messages(source): + """Every fprintf(stderr, ...) format string in `source`, with adjacent + literals concatenated (still escaped -- the matched prefix contains no + escapes, so matching against the raw literal is faithful to what the + runtime prints before the %d substitution).""" + messages = [] + for match in _FPRINTF.finditer(source): + parts = re.findall(_STRING, match.group(1)) + messages.append("".join(part[1:-1] for part in parts)) + return messages + + +class RefusalShapeCanaryTest(unittest.TestCase): + maxDiff = None + + @classmethod + def setUpClass(cls): + cls.messages = stderr_messages(COLIBRI_C.read_text(encoding="utf-8")) + + def messages_of(self, function): + return [m for m in self.messages if m.startswith(function + ":")] + + def test_each_function_still_names_a_refusal_the_matcher_greps(self): + for function in FUNCTIONS: + named = self.messages_of(function) + self.assertTrue( + named, + f"colibri.c has no stderr message starting '{function}:' -- the " + f"refusal site was renamed or removed, and the D-I2 bite matcher " + f"(test_fp8_serve_batch_e2e.REFUSAL) can no longer identify this " + f"death; update the matcher and the fleet bite lane together") + refusals = [m for m in named if "refus" in m] + self.assertTrue( + refusals, + f"{function}: no refusal-class stderr message left (the 'refus' " + f"discipline tests/test_qt_addrow.c pins at runtime) -- if the " + f"refusal moved, this canary and the e2e matcher must follow it") + self.assertTrue( + any(REFUSAL.search(m) for m in refusals), + f"none of {function}'s refusal messages matches the e2e " + f"REFUSAL regex {REFUSAL.pattern!r} -- an innocent reword " + f"silently turns the fleet bite proof's machine-match into a " + f"generic 'engine died'; keep one matchable refusal per " + f"function or change test_fp8_serve_batch_e2e.REFUSAL in the " + f"same commit. (Existential on purpose: an ADDITIONAL, " + f"differently-worded guard in this function is new coverage, " + f"not drift.)\nmessages: {refusals}") + + def test_matcher_stays_selective(self): + """The other drift direction: a loosened REFUSAL regex would make + expect-bite accept ANY death as the bite. Pin that it rejects a + representative non-refusal death line and the empty string, and + that it keys on the FUNCTION FAMILY, not the bare `unsupported + fmt=` words -- a matcher loosened to the words alone would accept + any future non-absorb `unsupported fmt=` message as the bite.""" + self.assertIsNone(REFUSAL.search("malloc: out of memory allocating 42 GB")) + self.assertIsNone(REFUSAL.search("colibri engine exited unexpectedly")) + self.assertIsNone(REFUSAL.search("")) + self.assertIsNone(REFUSAL.search("qt_resolve_fmt: unsupported fmt=9 refusing")) + # The REAL in-tree near-miss (not synthetic): layer_cuda_shard_kvb's + # refusal shares the `unsupported ... fmt=` words but is a different + # guard family (shard-layout, pinned by tests/test_shard_kvb_refuse.c) + # -- a loosened matcher that swallowed it would misfile a shard + # refusal as the absorb bite in a fleet lane. + self.assertIsNone(REFUSAL.search( + "layer_cuda_shard_kvb: unsupported kv_b fmt=5 for the head-shard " + "upload (only fmt 1/2/3/4 match the per-row byte/scale strides " + "computed here) -- refusing the shard")) + + +if __name__ == "__main__": + unittest.main() diff --git a/c/tests/test_fp8_serve_batch_e2e.py b/c/tests/test_fp8_serve_batch_e2e.py new file mode 100644 index 000000000..52bb35b62 --- /dev/null +++ b/c/tests/test_fp8_serve_batch_e2e.py @@ -0,0 +1,668 @@ +"""Defect-closure pin for the batched-serve fmt=8 abort (D-I2): two-slot +SERVE_BATCH mux against a REAL fp8 (fmt=8) `kv_b_proj` container. + +The defect: batched serve mode (openai_server's Engine always launches the +engine with SERVE_BATCH=1; --kv-slots > 1 gives the mux real slots) decodes +kv_b through the MLA absorb path -- qt_addrow / qt_matvec_rows in colibri.c. +Before the fmt=8 absorb branches landed, those functions had no fp8 case: +the first generation request against a container whose kv_b_proj resolves to +fmt=8 (fp8-e4m3-b128 -- both production fp8 container classes, f8_full and +f8x4g64, are in this class) killed the engine with the named refusal + + qt_addrow: unsupported fmt=8 for the per-row-scale absorb path ... + -- refusing rather than misread t->s[row]/t->q4 + +and the client saw a 500 `engine_error`. Single-stream chat never hit it, +so nothing in the FakeEngine-backed server suite could: the abort lives +below the engine wire protocol, on a code path only a real fmt=8 container +reaches. This test pins the closure end to end: serve the real container +batched, occupy BOTH KV slots concurrently, and require 200 + text on both +requests with the engine still alive afterwards. + +A real fmt=8 container is hundreds of GB and exists only on the model +hosts, never in CI, so discovery is by environment variable: + + COLI_FP8_CONTAINER root of an engine-loadable container whose + kv_b_proj is fp8 (fmt=8). Unset => named SKIP. + COLI_ENGINE engine binary override (same convention as coli); + default: the built `colibri` next to the server. + SET but wrong => loud FAIL, never a skip. + COLI_FP8_EXPECT_BITE=1 proof-of-bite mode for the fleet's old-binary + half: the run PASSES only if the invocation FAILS + to serve AND the captured server stderr carries + the named `unsupported fmt=` refusal. A death + without the refusal, or a clean serve, FAILS. + COLI_FP8_READY_TIMEOUT seconds to wait for the server to come up + (default 1800 -- container loads are large). + COLI_FP8_GENERATE_TIMEOUT seconds per generation request (default + 1800; two concurrent cold prefills on a + storage-bound host can be slow). + +Set-but-wrong is a FAILURE, not a skip: if COLI_FP8_CONTAINER names a +missing path, a container whose kv_b_proj is NOT fmt=8 (e.g. an int8 +spine), or a shard set with duplicate tensor names, this test would +otherwise pass (or fail for the wrong reason) without ever entering the +absorb branch it exists to pin -- the vacuous-gate hazard. The safetensors +headers are checked first (pure header reads, no tensor data) and the +mismatch is a loud failure naming what was found. The same doctrine covers +COLI_ENGINE: a set-but-nonexistent engine path fails, it does not skip. +(Known unreachable corner: a shape with O == ceil(O/128)*ceil(I/128) makes +per-row and per-block scale counts collide and the unstamped loader would +resolve fmt=1 where this check says fmt=8 -- real GLM kv_b O is in the +thousands, so no planned container can reach it.) + +The engine/server child environment is constructed fresh: ambient COLI_* +and engine-knob variables (KV8, KV_TQ, ...) are stripped except an +explicit backend/location allowlist plus two value-restricted knobs -- +exactly ABSORB=0 (the ratified non-absorb fleet arm; any other ABSORB +value is stripped), and exactly COLI_CUDA_ATTN=1 WHEN the lane opts in +with COLI_FP8_E2E_CUDA_ABSORB=1 (default: stripped, so a leaked ambient +COLI_CUDA_ATTN can never swap the code path under test) -- and a leaked +COLI_API_KEY cannot 401 the ready poll; the kept/dropped set is printed +and attached to every failure so the run artifact shows the env the +invocation actually saw. + +Lane semantics, stated exactly: by DEFAULT every lane (CPU or CUDA host) +exercises the CPU absorb arms -- the CUDA fmt=8 absorb decode is gated +on COLI_CUDA_ATTN (path selection) AND CUDA_DENSE (kv_b's +cuda_eligible), both of which this test strips unless the lane +explicitly opts in. A CUDA lane that sets COLI_FP8_E2E_CUDA_ABSORB=1 +(plus its COLI_CUDA/COLI_GPU bindings and ambient COLI_CUDA_ATTN=1 and +CUDA_DENSE=1) pins the CUDA absorb decode end-to-end, and the test then +REQUIRES the engine's GPU-dense boot line as a witness -- an opt-in run +whose engine reports resident-dense-on-CPU fails rather than banking a +vacuous green. Without the opt-in, the CUDA absorb arm is pinned only +by tests/test_backend_cuda.cu. +""" +import json +import os +import re +import signal +import socket +import struct +import subprocess +import sys +import threading +import time +import unittest +from pathlib import Path +from urllib.error import HTTPError, URLError +from urllib.request import Request, urlopen + +HERE = Path(__file__).resolve().parent +C_DIR = HERE.parent +CONTAINER = os.environ.get("COLI_FP8_CONTAINER") +ENGINE_OVERRIDE = os.environ.get("COLI_ENGINE") +EXPECT_BITE = os.environ.get("COLI_FP8_EXPECT_BITE") == "1" +READY_TIMEOUT = float(os.environ.get("COLI_FP8_READY_TIMEOUT", "1800")) +GENERATE_TIMEOUT = float(os.environ.get("COLI_FP8_GENERATE_TIMEOUT", "1800")) +PROMPT = "The primary colors are" # the DR-10 fixed prompt: short, deterministic continuation +MAX_TOKENS = 24 +# The refusal family this test exists to keep dead. Both absorb entry points +# carry the same named message shape; match the family, not one function, so +# a regression through either surfaces by name in the failure output. +REFUSAL = re.compile(r"(qt_addrow|qt_matvec_rows): unsupported fmt=") +ENGINE_DEATH = "colibri engine exited unexpectedly" + +# Child-environment policy (the invocation of record must not ride ambient +# state): backend selection and read-only model LOCATION are the only knob +# namespaces a lane may pass through -- CUDA lanes need COLI_CUDA/ +# COLI_GPU(S) per the DR-10 bindings, and a split/mirrored container needs +# the engine to be TOLD where its shards live (location config says where +# the weights ARE; a behavior knob changes what the engine DOES with them +# -- only the former class passes): +# COLI_MODEL_DIRS -- extra shard directories (st_init_multi's SPLIT +# layout); stripping it hides the operator's shards +# and fails as a bogus "not an fmt=8 container". +# COLI_MODEL_MIRROR -- read-only replica dirs (multi-SSD read fan-out). +# COLI_MMAP -- how weights are mapped from disk; load +# placement/feasibility, not decode semantics. +# COLI_DISK_WEIGHTS -- disk-resident weight policy; feasibility for +# hundreds-of-GB loads, not decode semantics. +# Every other COLI_* (COLI_API_KEY -> 401'd ready poll) and the bare +# engine knobs below (KV8/KV_TQ change the KV format of record; +# SNAP/SERVE/SERVE_BATCH/NGEN/KV_SLOTS belong to the server, which sets +# its own) are stripped, and the strip is recorded. Documented legacy +# aliases are stripped ALONGSIDE their primaries so a scrubbed primary +# cannot resurface under its old name: SNAP_MIRROR (consulted only when +# COLI_MODEL_MIRROR is unset/empty -- i.e. precisely after this scrub) is +# name-stripped, and TEMP is stripped only when FULLY NUMERIC (matching +# temp_from_env's strtod whole-string test: a numeric TEMP is the +# deprecated sampling alias, a path TEMP is the Windows/ROCm temp +# directory and must survive). Value-restricted exceptions: the ratified +# non-absorb fleet arm selects itself with ABSORB=0, so exactly that +# value passes (any other ambient ABSORB is stripped -- both decode paths +# at ABSORB=0/default are ratified-equivalent, so a leaked "0" can shift +# which path is pinned but never fake a pass). +# +# The CUDA-absorb opt-in (COLI_FP8_E2E_CUDA_ABSORB=1) is handled by +# INJECTION, not passthrough -- see CUDA_ABSORB_INJECT below. A first +# attempt value-allowed COLI_CUDA_ATTN=1 and CUDA_DENSE=1 through the scrub +# and relied on the LANE to set both ambient; the fresh-env allowlist only +# ever KEEPS a variable already in the environment, so when the lane set +# COLI_CUDA_ATTN=1 but not the bare (non-COLI) CUDA_DENSE, the eligibility +# knob never reached the child and the run booted resident-dense-on-CPU +# (the boot-line witness caught it as a hard failure). The opt-in is the +# single source of truth for the bundle the path needs, so the test now +# SETS it directly and reports it, instead of hoping the lane assembled +# the same set correctly. +ENV_KNOBS_ALLOWED = ("COLI_CUDA", "COLI_GPU", "COLI_GPUS", "COLI_NO_OMP_TUNE", + "COLI_MODEL_DIRS", "COLI_MODEL_MIRROR", "COLI_MMAP", + "COLI_DISK_WEIGHTS") +ENV_KNOBS_VALUE_ALLOWED = {"ABSORB": ("0",)} +ENV_KNOBS_STRIPPED = ("KV8", "KV_TQ", "ABSORB", "CAP_RAISE", "CUDA_DENSE", + "COLI_CUDA_ATTN", "SNAP", "SNAP_MIRROR", "SERVE", + "SERVE_BATCH", "NGEN", "KV_SLOTS") +CUDA_ABSORB_OPT_IN = "COLI_FP8_E2E_CUDA_ABSORB" +# The exact engine knobs the CUDA fmt=8 absorb decode needs, INJECTED into +# the child (with the lane's own COLI_CUDA=1 backend binding) when the lane +# opts in. COLI_CUDA_ATTN=1 selects the CUDA absorb dispatch; CUDA_DENSE=1 +# makes kv_b/o cuda_eligible -- qt_load grants eligibility only under +# g_cuda_dense (colibri.c:2123) and there is NO VRAM/budget gate on it +# (colibri.c:10966), so CUDA_DENSE=1 reaching the child is sufficient on a +# device that fits the dense tensors. CUDA_DENSE is a bare name (no COLI_ +# prefix), which is exactly why passthrough could not carry it and explicit +# injection is required. Both are stripped from ambient (above) so an +# unrequested value can never leak in; under the opt-in the test's own "1" +# is what the child sees, recorded in the kept-knobs line as injected. +CUDA_ABSORB_INJECT = {"COLI_CUDA_ATTN": "1", "CUDA_DENSE": "1"} + +# Boot-line witness for the opt-in arm: requesting a path is not proof it +# ran. The engine announces its residency decision at boot +# (colibri.c:11047-11049); under the opt-in this test asserts the GPU-dense +# line and REFUSES the CPU-dense line, so a misconfigured lane (e.g. the +# eligibility knob lost again) can never bank a vacuous green. +CUDA_DENSE_BOOT = "[CUDA] mode: routed experts + resident dense tensors" +CUDA_CPU_DENSE_BOOT = "[CUDA] mode: routed experts only (resident dense on CPU)" + + +def observed_cuda_boot_mode(stderr_text): + """The engine's `[CUDA] mode: ...` boot line as it appears in stderr, or + None if it never printed. Echoed by the witness so a PASS artifact + SHOWS the observed residency, not just an assertion that succeeded.""" + marker = "[CUDA] mode:" + start = stderr_text.find(marker) + if start == -1: + return None + end = stderr_text.find("\n", start) + return stderr_text[start:end if end != -1 else None] + + +def cuda_absorb_witness_failure(stderr_text): + """None when the boot-line witness holds for the CUDA-absorb opt-in; + otherwise the failure message. + + It checks two things against the engine's stderr: that the + resident-dense-on-CPU boot line is absent, and that the + routed+resident-dense line is present. Module-level so it can be called + without constructing the test case, and so it can be exercised directly + once there is a caller outside the container-gated test below -- there is + none today, so on a machine without a container this function has no + coverage at all.""" + if CUDA_CPU_DENSE_BOOT in stderr_text: + return ("CUDA-absorb opt-in was requested but the engine booted with " + f"{CUDA_CPU_DENSE_BOOT!r} -- resident dense stayed on the CPU, " + "kv_b never became cuda_eligible, and the CUDA absorb path " + "cannot have run (vacuous green refused; check CUDA_DENSE=1 " + "reached the child, see the kept-knobs line)") + if CUDA_DENSE_BOOT not in stderr_text: + return ("CUDA-absorb opt-in was requested but the engine's boot " + f"witness {CUDA_DENSE_BOOT!r} never appeared on stderr -- " + "cannot certify the CUDA absorb path ran") + return None + + +def _numeric_temp(value): + """temp_from_env's whole-string strtod test: TEMP acts as the deprecated + sampling alias only when fully numeric; a path value is the system temp + directory and is not a knob.""" + if not value: + return False + try: + float(value) + return True + except ValueError: + return False + + +def _child_env(): + """Fresh server/engine environment per the policy above. Returns + (env, report): the report names kept knobs with values and dropped + knobs by NAME only (a dropped COLI_API_KEY must not leak its value + into a run artifact).""" + opt_in = os.environ.get(CUDA_ABSORB_OPT_IN) == "1" + env, kept, dropped = {}, [], [] + for name, value in os.environ.items(): + # The CUDA-absorb bundle is INJECTED below when opted in, never taken + # from ambient -- so the child sees the test's "1", not a leaked + # value, and an unrequested ambient copy is always stripped here. + if name in CUDA_ABSORB_INJECT: + dropped.append(name) + elif name in ENV_KNOBS_ALLOWED or value in ENV_KNOBS_VALUE_ALLOWED.get(name, ()): + env[name] = value + kept.append(f"{name}={value}") + elif name.startswith("COLI_") or name in ENV_KNOBS_STRIPPED: + dropped.append(name) + elif name == "TEMP" and _numeric_temp(value): + dropped.append(name) # numeric TEMP = deprecated COLI_TEMP alias + else: + env[name] = value # non-knob system env (PATH, HOME, LD_LIBRARY_PATH, ...) + if opt_in: + for name, value in CUDA_ABSORB_INJECT.items(): + env[name] = value + kept.append(f"{name}={value} (injected)") + dropped = [d for d in dropped if d not in CUDA_ABSORB_INJECT] + return env, ("kept knobs: " + (", ".join(sorted(kept)) or "(none)") + + "; dropped knobs: " + (", ".join(sorted(dropped)) or "(none)")) + + +def _default_engine(): + """The built glm engine next to openai_server.py -- same candidates, + same order, as the server's own default_engine().""" + for name in ("colibri", "glm"): + for suffix in ("", ".exe"): + candidate = C_DIR / (name + suffix) + if candidate.exists(): + return candidate + return C_DIR / "colibri" + + +ENGINE = Path(ENGINE_OVERRIDE) if ENGINE_OVERRIDE else _default_engine() + + +def _shard_headers(container): + """Merged {tensor name: (shape, nbytes)} across every *.safetensors shard + header in the container root, plus the duplicate-name collisions found on + the way. Header-only reads (u64 length + JSON): no tensor data is + touched, so scanning a 140-shard container is cheap. Duplicates are + returned rather than silently last-shard-wins: the engine refuses a + duplicate tensor name across shards (st_init's D-2 guard), so a merged + map that picked either copy could disagree with the load in either + direction.""" + headers, owner, duplicates = {}, {}, [] + shards = sorted(Path(container).glob("*.safetensors")) + for shard in shards: + with open(shard, "rb") as f: + (hlen,) = struct.unpack(" 0, + f"COLI_FP8_CONTAINER={CONTAINER} holds no *.safetensors shards") + self.assertFalse( + duplicates, + f"COLI_FP8_CONTAINER={CONTAINER} carries duplicate tensor names " + f"across shards (the engine's st_init refuses exactly this, and " + f"this check could otherwise judge a shard the engine never " + f"loads): {'; '.join(duplicates[:5])}") + is_fmt8, detail = _kv_b_fmt8_check(headers) + self.assertTrue(is_fmt8, + f"COLI_FP8_CONTAINER={CONTAINER} is not an fmt=8 kv_b " + f"container; this test would be vacuous against it: {detail}") + + server = _Server(_free_port()) + self.addCleanup(server.close) + + # Startup: poll /v1/models (touches only the server, not the engine + # generate path) until it answers, and read the served model id from + # it rather than hardcoding one. Sleep EVERY iteration -- a fast + # non-200 answer must not turn the poll into a hot spin -- and fail + # fast on auth/host-guard rejections, which no amount of waiting + # will turn into a 200. + deadline = time.time() + READY_TIMEOUT + model_id = None + while time.time() < deadline: + if server.proc.poll() is not None: + self._fail_with_server_evidence( + server, f"server exited rc={server.proc.returncode} before READY") + status = None + try: + status, body = _get(server.url("/v1/models"), timeout=10) + except (URLError, OSError): + pass # not accepting yet -- keep waiting + if status == 200: + model_id = json.loads(body)["data"][0]["id"] + break + if status in (401, 403): + self._fail_with_server_evidence( + server, f"ready poll got HTTP {status} from /v1/models -- " + f"an auth/host-guard rejection, not a slow load; " + f"waiting longer cannot fix it") + time.sleep(2) + if model_id is None: + self._fail_with_server_evidence( + server, f"server not ready within {READY_TIMEOUT:.0f}s") + + # The D-I2 invocation class: two generation requests IN FLIGHT + # TOGETHER, pinned to distinct KV slots so the 2-slot mux really + # multiplexes (conversation hashing would put one identical prompt in + # one slot). temperature 0 / fixed prompt per the invocation of + # record. On a pre-fix engine the FIRST decode through the absorb + # path kills the engine and one or both of these come back 500. + results = [None, None] + + def ask(slot): + results[slot] = _post(server.url("/v1/completions"), { + "model": model_id, "prompt": PROMPT, "max_tokens": MAX_TOKENS, + "temperature": 0, "cache_slot": slot, + }, timeout=GENERATE_TIMEOUT) + + threads = [threading.Thread(target=ask, args=(slot,)) for slot in (0, 1)] + for t in threads: + t.start() + for t in threads: + t.join(GENERATE_TIMEOUT + 60) + + if EXPECT_BITE: + self._assert_bite(server, results) + return + + for slot, outcome in enumerate(results): + self.assertIsNotNone(outcome, f"slot {slot}: request never completed " + f"(asking thread still blocked)") + kind = outcome[0] + if kind == "timeout": + self._fail_with_server_evidence( + server, f"slot {slot}: generation request timed out after " + f"{outcome[1]:.0f}s (COLI_FP8_GENERATE_TIMEOUT to raise; " + f"a timeout is not a hang and not a protocol failure)") + if kind == "neterr": + self._fail_with_server_evidence( + server, f"slot {slot}: connection failed mid-request " + f"({outcome[1]}) -- engine/server dropped the stream") + _, status, body, resp_headers = outcome + if status != 200: + self._fail_with_server_evidence( + server, + f"slot {slot}: HTTP {status} (the D-I2 defect signature is a " + f"500 engine_error here)\nresponse body: " + f"{body.decode(errors='replace')[:2000]}") + payload = json.loads(body) + text = payload["choices"][0]["text"] + # "Coherent content", pinned mechanically: nonempty text with at + # least one real word, and the engine accounted for generated + # tokens. (Semantic quality belongs to the numeric battery, not + # this defect pin.) + self.assertTrue(text.strip(), + f"slot {slot}: 200 with empty text -- content-free " + f"success is not closure") + self.assertRegex(text, r"[A-Za-z]{2}", + f"slot {slot}: no word-like content in {text!r}") + self.assertGreater(payload["usage"]["completion_tokens"], 0, + f"slot {slot}: usage reports zero generated tokens") + # Concurrency witness: with two requests pinned to two distinct + # free slots, neither should queue. A silently serialized mux + # would park the second request for the first one's whole + # generation -- minutes, not milliseconds -- while every other + # assertion here stayed green. + queue_wait = resp_headers.get("x-colibri-queue-wait-ms") + self.assertIsNotNone(queue_wait, + f"slot {slot}: no x-colibri-queue-wait-ms response " + f"header -- the admission witness is gone") + self.assertLess(float(queue_wait), 1000.0, + f"slot {slot}: queued {queue_wait} ms behind the other " + f"slot -- the 2-slot mux is serializing, not batching") + # Echo the observed admission latency: the concurrency witness is + # a positive coverage claim, so its measured value belongs in the + # PASS artifact, not only in the failure path. + print(f"[fp8-e2e] witness: slot {slot} admission queue-wait = " + f"{queue_wait} ms", flush=True) + + # No engine death, three ways: the named refusal family never fired, + # the server never recorded an engine exit, and the server still + # answers after both generations. + stderr = server.stderr() + refusal = REFUSAL.search(stderr) + self.assertIsNone( + refusal, "the engine printed the absorb-path refusal this test " + f"exists to keep dead:\n{stderr[stderr.rfind('qt_'):][:500]}") + self.assertNotIn(ENGINE_DEATH, stderr, + "the engine died during the batched run") + self.assertIsNone(server.proc.poll(), "server process exited mid-test") + status, _ = _get(server.url("/v1/models"), timeout=30) + self.assertEqual(status, 200, "server stopped answering after the batched pair") + + # Boot-line witness (opt-in arm only): requesting the CUDA absorb + # path is not proof it ran -- certify it from the engine's own boot + # report, and refuse the CPU-dense line outright. The observed line + # is ECHOED on both pass and fail (positive coverage must be a + # captured artifact, not an inference from a green assertion), so a + # CI log carries the routed+dense observation verbatim. + if os.environ.get(CUDA_ABSORB_OPT_IN) == "1": + observed = observed_cuda_boot_mode(stderr) + print(f"[fp8-e2e] witness: observed boot mode = {observed!r}", + flush=True) + witness_failure = cuda_absorb_witness_failure(stderr) + if witness_failure: + self._fail_with_server_evidence(server, witness_failure) + + +if __name__ == "__main__": + unittest.main() diff --git a/c/tools/repack_fp8_passthrough.py b/c/tools/repack_fp8_passthrough.py index 3e4d656f6..c54528980 100644 --- a/c/tools/repack_fp8_passthrough.py +++ b/c/tools/repack_fp8_passthrough.py @@ -650,7 +650,7 @@ def _check_stamp_budget(total_stamped): # produce ONE container, but each keeps its own progress manifest (a resume # manifest has to be per-selection -- see --mtp's comment in main()). The stamp # budget is not per-selection: st_fmt_stamp_ingest accumulates S->fmt_n over every -# shard it discovers in the directory (c/st.h:372), so ST_FMT_STAMP_MAX is a +# shard it discovers in the directory (c/st.h, the stamp scan), so ST_FMT_STAMP_MAX is a # CONTAINER-wide bound. Counted per-manifest, two passes could each stay under the # cap and still hand the engine a container over it -- the writer guarantee would # be 2x the reader's bound. So the budget check sums this pass's running total with @@ -926,17 +926,19 @@ def _print_inventory_summary(all_inv, dry_run): # the minted directory does not have). # # The EXTERNAL (non-tensor) files the loader/server actually open from a -# model dir at runtime, verified against this worktree's HEAD: -# - config.json cfg_root, colibri.c:1361 -- fopen(...); if(!f){ +# model dir at runtime. Cited by SYMBOL, not line number: the previous table +# carried line anchors that rotted twice, and a stale number is worse than no +# number because it reads as a verification that was not performed. +# - config.json colibri.c, cfg_root() -- fopen(...); if(!f){ # perror(p); exit(1); } -- MANDATORY, the run aborts without it. Also -# read by openai_server.py's Engine.__init__ (:1773) for arch detection. -# - generation_config.json colibri.c:1404-1405 -- fopen, comment "assente +# read by openai_server.py's Engine.__init__ (:2635) for arch detection. +# - generation_config.json colibri.c, cfg_root() -- fopen, comment "assente # = nessun problema: e' opzionale" -- best-effort; HF's authority for # generation defaults (extra EOS stop ids) when present. -# - tokenizer.json c/tok.h:101 (tk_read_file, called from tok_load) +# - tokenizer.json c/tok.h, tk_read_file() (called from tok_load) # -- fopen(...); if(!f){ perror(path); exit(1); } -- MANDATORY. Called # from every serve/generate entry point that needs a tokenizer -# (colibri.c:7294 run_text, :7952 run_serve_mux, :8134 main serve loop). +# (colibri.c: run_text, run_serve_mux, and the main serve loop). # Before this fix main() never copied it into --outdir, so a minted # directory was not standalone-loadable for THIS reason alone (confirmed # by V2's end-to-end smoke test, 2026-08-18: load only succeeded via a @@ -947,8 +949,8 @@ def _print_inventory_summary(all_inv, dry_run): # fopen/Path().open against a model dir) and therefore excluded: # - tokenizer_config.json, chat_template.jinja -- the GLM-5.2 chat template # is reimplemented natively in code, not read from the .jinja file -# (colibri.c:8211 comment: "template UFFICIALE GLM-5.2 (chat_template -# .jinja): niente \n dopo i ruoli..."; openai_server.py:1016: "AUTHORITATIVE +# (colibri.c, comment: "template UFFICIALE GLM-5.2 (chat_template +# .jinja): niente \n dopo i ruoli..."; openai_server.py: "AUTHORITATIVE # GLM-5.2 tool-declaration block (byte-matches chat_template.jinja)" -- # both hardcoded to match the file's behavior, not sourced from it). # - README.md, LICENSE, .gitattributes -- pure repo/documentation metadata, diff --git a/docs/FORMATS.md b/docs/FORMATS.md index b52e90c8d..16f38914a 100644 --- a/docs/FORMATS.md +++ b/docs/FORMATS.md @@ -101,8 +101,9 @@ allowlist, :3379), consulted by both gate sites — `attention_rows` (:3448) and `layer_forward_rows` (:5792) — and by the load-time notice `metal_fmt_gate_notice` (:1866, called from `model_init`). -Sources for all rows (`c/quant.h`/`c/colibri.c` line numbers at this PR -pair's current restack, base dev `292ed4c`): +Sources for all rows (`c/quant.h`/`c/colibri.c` line numbers verified at +this branch's own head -- originally written against base dev `292ed4c`, +re-anchored here because line numbers rot with the file, not the base): - **fmt=0/1/2/3** — allocation policy: `qt_alloc`, `c/colibri.c` (`bits>=16→fmt=0`, `bits>=5→fmt=1`, `bits>=4→fmt=2`, else `fmt=3`). From b95b942197f585bfc191d204167e222b26c005c9 Mon Sep 17 00:00:00 2001 From: monotophic Date: Tue, 22 Sep 2026 22:58:48 -0400 Subject: [PATCH 5/6] docs(fp8): correct stale anchors, an unreachable test rule and a wrong test comment Five accuracy items from the r7 audit (AUDIT_1102_r7_ea8a7e4d.md, findings F1-F5), none touching the decode path. docs/FORMATS.md's Metal fused-gate paragraph (F1) still named a colibri.c line number for each of five symbols and cited a non-upstream branch (kvb/fmt-gate-notice-r4) as its source of truth, even though the "Sources for all rows" table just below it had already been de-anchored. All five are now symbol references in the same style as that table, and the branch name is gone. repack_fp8_passthrough.py's EXTERNAL-files table (F2) de-anchored its line numbers everywhere the preamble above it promises, except one: the Engine.__init__ reference still carried a line number (previously wrong per r6 F4, and edited to a different wrong number in r7). It is now a bare symbol reference, matching the rest of the block. tests/test_e8x4g64_loader$(EXE) (F3) had a rule and a comment claiming it compiles under the suite's own CFLAGS so its warnings are caught, but the rule was unreachable from every target -- test-c, test, check, and all -- since TEST_EXCLUDE drops it from TEST_BINS and nothing else names it as a prerequisite. Added e8x4g64-loader-check, a phony build-only target shaped like qwen38-tier-engine-check, so `make e8x4g64-loader-check` actually builds it under those flags on demand. check itself is untouched: the new target is not a prerequisite of check/test/test-c/all, so this adds no time to `make check`. test_qt_addrow.c's block-scale comment (F4) said a ZERO fmt=8 scale is "refused by name" alongside NaN. It is not: colibri.c is explicit that a zero scale is valid and must decode to zeros, and test_fmt8_zero_scale_decodes_to_zeros twenty lines below already asserts exactly that. The comment now describes only the NaN guard the two calls below it exercise, and points at that assertion for the zero case. fp8_format.h (F5) is added beside quant.h/backend_cuda.cu on two of the three Makefile rules the audit named: tests/test_kimi_cuda_expert$(EXE) (kimi_k3.c includes quant.h directly, which pulls in fp8_format.h) and tests/bench_cuda_resident_batch$(EXE) (backend_cuda.cu -- a separate translation unit built into the same binary -- includes fp8_format.h directly, a miss no quant.h-based framing could see). The third, tests/test_qwen36_dnproj_batch$(EXE), is deliberately left alone: full preprocessing under both its build variants (plain and the -DCOLI_DNPROJ_REAL_CUDA alternate) confirms qwen36.c never reaches quant.h or fp8_format.h at all, exactly as 6826c00a's own commit message already documented when it introduced fp8_format.h. The audit's framing of this rule as a regression does not hold up under that check; the pre-existing quant.h entry on that rule is itself already a no-op and is not this change's to fix. Not in scope: the fmt=8 NaN exit(1)-inside-omp policy, disclosed in the r7 thread note and left for a later round. --- c/Makefile | 16 +++++++++++++--- c/tests/test_qt_addrow.c | 10 +++++----- c/tools/repack_fp8_passthrough.py | 2 +- docs/FORMATS.md | 11 +++++------ 4 files changed, 24 insertions(+), 15 deletions(-) diff --git a/c/Makefile b/c/Makefile index eed827041..749737ca6 100644 --- a/c/Makefile +++ b/c/Makefile @@ -501,7 +501,9 @@ TEST_RULES := $(shell sed -n 's|^tests/\(test_[a-z0-9_]*\)\$$(EXE):.*|\1|p' $(f # test_e8x4g64_loader takes a minted container directory on argv and prints its # usage (exit 2) without one, so it is a harness rather than a gate: it has a # rule here so it compiles under the suite's own $(CFLAGS) and its warnings are -# caught, and tests/test_e8x4g64_mint_load.py is what actually drives it. +# caught -- reachable via `make e8x4g64-loader-check`, kept out of `check` itself +# since a bare invocation exits 2 -- and tests/test_e8x4g64_mint_load.py is what +# actually drives it end to end. # test_qwen38_tier_engine drives qwen38.c over the generated FP8 fixture # (qwen38_tiny_fp8, gitignored); it runs from qwen38-tier-engine-check. TEST_EXCLUDE = test_uring test_deepseek_v4 test_v4_ownership test_v4_serve_framing \ @@ -1087,7 +1089,7 @@ fp8-bench: fp8_bench$(EXE) ./fp8_bench$(EXE) # Matched resident int8 GPU projection timings; run explicitly, not in CI. -tests/bench_cuda_resident_batch$(EXE): tests/bench_cuda_resident_batch.cu backend_cuda.cu backend_cuda.h backend_gpu_compat.h +tests/bench_cuda_resident_batch$(EXE): tests/bench_cuda_resident_batch.cu backend_cuda.cu backend_cuda.h fp8_format.h backend_gpu_compat.h "$(GPUCC)" $(GPUFLAGS) backend_cuda.cu tests/bench_cuda_resident_batch.cu -o $@ $(ANS_NVCC_LIBS) $(GPU_TEST_LIBS) # NOCUDA_*: olmoe.c has no COLI_CUDA code at all, so CUDA=1 would otherwise @@ -2109,6 +2111,14 @@ tests/test_shard_kvb_refuse$(EXE): tests/test_shard_kvb_refuse.c colibri.c st.h tests/test_e8x4g64_loader$(EXE): tests/test_e8x4g64_loader.c st.h quant.h fp8_format.h compat.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) +# Reachable build-only entry point for the harness above: `make check` never +# calls this (it stays out of TEST_BINS/test-c, per TEST_EXCLUDE), so this adds +# no time to `check`. It exists so the suite's own $(CFLAGS) and warnings are +# actually exercised on demand, the way qwen38-tier-engine-check does for +# test_qwen38_tier_engine below. +.PHONY: e8x4g64-loader-check +e8x4g64-loader-check: tests/test_e8x4g64_loader$(EXE) + test-c: $(TEST_BINS) $(PYTHON) tools/run_tests.py $(TEST_BINS) @@ -2247,5 +2257,5 @@ tests/test_kimi_request_state$(EXE): tests/test_kimi_request_state.c kimi_k3.c k $(CC) $(NOCUDA_CFLAGS) tests/test_kimi_request_state.c $(VK_OBJ) -o tests/test_kimi_request_state$(EXE) $(NOCUDA_LDFLAGS) # Kimi CUDA dispatch/fallback with a fake backend; no CUDA toolkit required. -tests/test_kimi_cuda_expert$(EXE): tests/test_kimi_cuda_expert.c kimi_k3.c backend_cuda.h kv_prefix.h serve_codec.h st.h json.h tok.h tok_unicode.h tok_unicode_o200k.h compat.h quant.h idot.h omp_tune.h route_trace.h +tests/test_kimi_cuda_expert$(EXE): tests/test_kimi_cuda_expert.c kimi_k3.c backend_cuda.h kv_prefix.h serve_codec.h st.h json.h tok.h tok_unicode.h tok_unicode_o200k.h compat.h quant.h fp8_format.h idot.h omp_tune.h route_trace.h $(CC) $(NOCUDA_CFLAGS) -DCOLI_CUDA $< $(VK_OBJ) -o $@ $(NOCUDA_LDFLAGS) diff --git a/c/tests/test_qt_addrow.c b/c/tests/test_qt_addrow.c index 254c9e772..6cac5986c 100644 --- a/c/tests/test_qt_addrow.c +++ b/c/tests/test_qt_addrow.c @@ -390,12 +390,12 @@ static int expect_refuse_call(absorb_fn fn, const char *tag){ #endif } -/* fmt=8 block scales that cannot be multiplied through. A NaN scale poisons - * every accumulator downstream of it; a ZERO scale emits an all-zero block, - * which is the same output a decode against an unwritten table produces -- so - * both are refused by name rather than propagated. One probe per value per +/* A NaN fmt=8 block scale cannot be multiplied through: it poisons every + * accumulator downstream of it, so it is refused by name rather than + * propagated. (A ZERO scale is the opposite case -- valid data, not refused; + * see test_fmt8_zero_scale_decodes_to_zeros below.) One probe per value per * function; the scale is poisoned AFTER fill_fmt8 has built a valid tensor, so - * the only thing under test is the new guard. */ + * the only thing under test is the NaN guard. */ static void poison_scale(QT *t, float v){ t->s[0]=v; } static void call_addrow_fmt8_nan(void){ diff --git a/c/tools/repack_fp8_passthrough.py b/c/tools/repack_fp8_passthrough.py index c54528980..71d097287 100644 --- a/c/tools/repack_fp8_passthrough.py +++ b/c/tools/repack_fp8_passthrough.py @@ -931,7 +931,7 @@ def _print_inventory_summary(all_inv, dry_run): # number because it reads as a verification that was not performed. # - config.json colibri.c, cfg_root() -- fopen(...); if(!f){ # perror(p); exit(1); } -- MANDATORY, the run aborts without it. Also -# read by openai_server.py's Engine.__init__ (:2635) for arch detection. +# read by openai_server.py's Engine.__init__ for arch detection. # - generation_config.json colibri.c, cfg_root() -- fopen, comment "assente # = nessun problema: e' opzionale" -- best-effort; HF's authority for # generation defaults (extra EOS stop ids) when present. diff --git a/docs/FORMATS.md b/docs/FORMATS.md index 16f38914a..c0d3c4899 100644 --- a/docs/FORMATS.md +++ b/docs/FORMATS.md @@ -94,12 +94,11 @@ fused path) — AND the other fused-bound tensors (`q_a`, `q_b`, `kv_a`, `o`, and on sparse layers `sh_gate`/`sh_up`/`sh_down`) sit on the fmt 1/2/3/4 allowlist; any other format on any of those tensors makes the affected layers' decode take the CPU path instead, announced by a one-line-per-tensor-kind -`[METAL]` stderr notice at load. Single source of truth (all anchors -`c/colibri.c` at branch head `kvb/fmt-gate-notice-r4`): the shared per-layer -predicate `metal_fused_layer_fmt_miss` (:3396, over the `metal_fused_fmt_ok` -allowlist, :3379), consulted by both gate sites — `attention_rows` (:3448) and -`layer_forward_rows` (:5792) — and by the load-time notice -`metal_fmt_gate_notice` (:1866, called from `model_init`). +`[METAL]` stderr notice at load. Single source of truth in `c/colibri.c`: the +shared per-layer predicate `metal_fused_layer_fmt_miss`, over the +`metal_fused_fmt_ok` allowlist, consulted by both gate sites — +`attention_rows` and `layer_forward_rows` — and by the load-time notice +`metal_fmt_gate_notice`, called from `model_init`. Sources for all rows (`c/quant.h`/`c/colibri.c` line numbers verified at this branch's own head -- originally written against base dev `292ed4c`, From 4159a23459ec65ec794b6c5888181fba91f95506 Mon Sep 17 00:00:00 2001 From: monotophic Date: Tue, 22 Sep 2026 23:15:52 -0400 Subject: [PATCH 6/6] docs(fp8): FORMATS.md names symbols, so stop promising line numbers --- docs/FORMATS.md | 19 ++++++++++--------- 1 file changed, 10 insertions(+), 9 deletions(-) diff --git a/docs/FORMATS.md b/docs/FORMATS.md index c0d3c4899..374741607 100644 --- a/docs/FORMATS.md +++ b/docs/FORMATS.md @@ -42,9 +42,9 @@ own verification anchor in its sources bullet): the stamp+registry series (#529) stacked directly on the fp8-passthrough series (#528), which is in turn based on dev `292ed4c` (post-#465, post-#457 Metal grouped-GEMV merge, post-#705 Vulkan/Kimi-K3 MXFP4 merge) — no -cross-tree line-number mixing. Every `c/colibri.c`/`c/quant.h` line number -in this document reflects that restack; re-verify them again if this branch -is rebased further. The fmt=6 and fmt=7 rows are upstream's own merged code +cross-tree mixing. Every `c/colibri.c`/`c/quant.h` symbol named in this +document is verified present at this branch's own head. The fmt=6 and +fmt=7 rows are upstream's own merged code (this branch's only fmt=6-adjacent change is the collision handling inside `qt_resolve_fmt`, `c/colibri.c`; it does not touch fmt=7/MXFP4 at all). @@ -100,9 +100,10 @@ shared per-layer predicate `metal_fused_layer_fmt_miss`, over the `attention_rows` and `layer_forward_rows` — and by the load-time notice `metal_fmt_gate_notice`, called from `model_init`. -Sources for all rows (`c/quant.h`/`c/colibri.c` line numbers verified at -this branch's own head -- originally written against base dev `292ed4c`, -re-anchored here because line numbers rot with the file, not the base): +Sources for all rows (`c/quant.h`/`c/colibri.c` symbols named below, +verified present at this branch's own head -- originally identified +against base dev `292ed4c`, reconfirmed here against the current +restack): - **fmt=0/1/2/3** — allocation policy: `qt_alloc`, `c/colibri.c` (`bits>=16→fmt=0`, `bits>=5→fmt=1`, `bits>=4→fmt=2`, else `fmt=3`). @@ -153,9 +154,9 @@ re-anchored here because line numbers rot with the file, not the base): override that default — see "The metadata stamp" below for the exact rule in both cases. FMT_NAMES table (`name string` to `fmt int`): `c/colibri.c`. -- **no ordinal** (`int4-rans256-g0`, merged tools-only tier — line numbers - at dev `7fb1159`, post-#671 merge `a3a5a75`, not at this PR pair's - restack base) — codec + record reader/writer: `c/rans.h` +- **no ordinal** (`int4-rans256-g0`, merged tools-only tier — symbols + verified at dev `7fb1159`, post-#671 merge `a3a5a75`, not at this PR + pair's restack base) — codec + record reader/writer: `c/rans.h` (`RANS_NSTREAMS 256`, `c/rans.h`; the record layout in the file-header comment, `c/rans.h`; that same header names its engine consumer "a future engine decode stage", `c/rans.h` — the format's own statement