Session 11 work. Saves the llama.cpp source modifications as out-of-tree patches at scripts/llama_cpp_patches/ so they survive git pull on llama.cpp (same approach as the existing slot-cap patch). Includes the updated SESSION_HANDOFF.md with full context for picking up the remaining FA work next session. What landed (in ~/llama.cpp/): Stage 5 — CLI flag wiring: - common/arg.cpp registers tq3_k256 as a --cache-type-k/v value plus a tq3 short alias. - llama-context.cpp gains a model-load-time pre-scan that warns when no layers match head_dim=256 (correctly handles ISWA models like Gemma 4 E4B which build two separate cache objects). - llama-kv-cache.cpp gains a per-layer match-branch log so every layer's effective cache type is observable, not just fallbacks. Bug fixes discovered during dispositive validation: - Bug ultraworkers#1 (Stage 5 ISWA scoping): warning was per-cache; moved to per-model in llama-context.cpp. Was actively misleading users on Gemma 4 E4B (the target model). - Bug ultraworkers#3 (Stage 4 missing match-branch log): only fallback layers logged. Now all layers log effective type. - Bug ultraworkers#4 (CPU type_traits_cpu missing TQ3_K256 entry): SIGSEGV in ggml_compute_forward_set_rows from NULL .from_float. Stage 1 only registered the type in ggml.c base traits, missing ggml-cpu/ggml-cpu.c's per-CPU table. Found via gdb. Added the entry plus the quantize_row_tq3_k256 wrapper in ggml-cpu/quants.c. - Bug ultraworkers#5 (CPU FA missing vec_dot for TQ3_K256): second SIGSEGV in ggml_compute_forward_flash_attn_ext from NULL .vec_dot. Added ggml_vec_dot_tq3_k256_f32 (slow reference: dequant block to stack buffer, f32 dot with Q) and wired vec_dot_type=F32. Stage 3 — CUDA SET_ROWS kernel: - turboquant.cu (257 lines): single self-contained TU with the per-block quantize device function, the SET_ROWS kernel, host launch wrappers (i64/i32 index variants), the device-side Pi matrix + Lloyd-Max codebook, and lazy host init via std::call_once. Single-TU design avoids the cross-TU __device__ symbol issue nvcc has without RDC. - turboquant.cuh: declares the host-callable entry points consumed by set-rows.cu. - set-rows.cu dispatch gains a TQ3_K256 branch. - ggml-cuda.cu adds TQ3_K256 to the GGML_OP_SET_ROWS allowlist. Validation status: - CPU path (-ngl 0 --cache-type-k tq3_k256 --cache-type-v tq3_k256): end-to-end coherent text on Gemma 4 E4B at ~1.7 tok/s. Per-layer log confirms 20 SWA tq3_k256 + 4 FA q4_0 fallback. Bug ultraworkers#1 warning correctly does NOT fire. - GPU SET_ROWS only: KV cache allocates on CUDA0 (was 0 MiB, now 7.66 MiB). Original ggml-backend.cpp:809 abort is gone. - GPU end-to-end: still blocked on CUDA Flash Attention dequant for tq3_k256. Same architectural shape as Bug ultraworkers#5 on the CPU side. Documented as Stage 3.5 in SESSION_HANDOFF.md with three implementation paths (vec template extension, custom kernel in turboquant.cu, naive thread-0 vec_dot) and the per-thread Q layout subtleties that make this multi-hour work. The patches directory has a README.md explaining apply order. Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>

这是好事啊 #3

Description

Metadata

Metadata

Assignees

Labels

Type

Fields

Projects

Milestone

Relationships

Development

Issue actions