diff --git a/c/Makefile b/c/Makefile index 61bc0a996..7f3775f23 100644 --- a/c/Makefile +++ b/c/Makefile @@ -498,12 +498,19 @@ 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 -- 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 \ 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) @@ -766,7 +773,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 exact_dot.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 exact_dot.h 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. @@ -785,7 +792,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). @@ -815,7 +822,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; } @@ -834,7 +841,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 $@ @@ -957,7 +964,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) @@ -1068,7 +1075,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) @@ -1076,14 +1083,14 @@ 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) ./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 @@ -1132,7 +1139,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) @@ -1140,7 +1147,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 @@ -1251,10 +1258,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 serve_budget.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 serve_budget.h 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 @@ -1282,7 +1289,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 @@ -1297,10 +1304,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 @@ -1311,19 +1318,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 @@ -1345,13 +1352,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 @@ -1397,12 +1404,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 @@ -1418,7 +1425,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 @@ -1453,13 +1460,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 serve_budget.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 serve_budget.h 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 @@ -1477,33 +1484,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 @@ -1512,15 +1519,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 @@ -1547,13 +1554,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 @@ -1562,10 +1569,17 @@ 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 +# 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_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_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 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 @@ -1583,7 +1597,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 @@ -1628,7 +1642,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 $@ @@ -1654,7 +1668,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 $@ @@ -1820,7 +1834,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 @@ -1949,7 +1963,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 @@ -1971,15 +1985,15 @@ tests/test_exact_dot$(EXE): tests/test_exact_dot.c exact_dot.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) @@ -1990,29 +2004,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. @@ -2022,34 +2036,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 @@ -2058,7 +2072,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 @@ -2097,6 +2111,29 @@ 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) + +# 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) + +# 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) @@ -2231,9 +2268,9 @@ 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. -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/Makefile.deepseek-v4 b/c/Makefile.deepseek-v4 index 7986c748e..97df35f70 100644 --- a/c/Makefile.deepseek-v4 +++ b/c/Makefile.deepseek-v4 @@ -264,20 +264,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 $(V4_FLAGS_STAMP) $(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 $(V4_FLAGS_STAMP) $(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 $(V4_FLAGS_STAMP) + tensor.h quant.h fp8_format.h native_quant.h native_quant_batch.h $(V4_FLAGS_STAMP) $(CC) $(CFLAGS) -DCOLI_V4_TEST_HOOKS \ -DCOLI_V4_UNIT_NATIVE_QUANT_BATCH -c deepseek_v4.c -o $@ 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/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 4f22a587b..185297ad6 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 @@ -2278,6 +2279,42 @@ 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 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 + * 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; @@ -3870,6 +3907,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 @@ -3903,20 +3981,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); } @@ -3966,18 +4040,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 new file mode 100644 index 000000000..d10ea461e --- /dev/null +++ b/c/fp8_format.h @@ -0,0 +1,35 @@ +/* fmt=8 (fp8-e4m3-b128) block geometry -- the single definition site for the + * 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, 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 + +#include + +#define FP8_BLOCK 128 + +/* 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/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_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_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_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_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/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/c/tests/test_qt_addrow.c b/c/tests/test_qt_addrow.c index 46c3673ca..6cac5986c 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;i +#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) 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 of this two-sided integration now exists on both backends: +colibri.c's MLA-absorption path (qt_addrow/qt_matvec_rows, called only on +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 @@ -657,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 @@ -933,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__ 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 @@ -954,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 c7724d8aa..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). @@ -94,39 +94,40 @@ 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`). - -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` +`[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` 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`). - 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 +140,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`. -- **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` + `c/colibri.c`. +- **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 + 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 +170,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