Research: port V100/SM70 kernel and runtime optimizations into llama.cpp-fork #9

Activity

1. novkien commented on Oct 1, 2026

   

   novkien

   on Oct 1, 2026 – with ChatGPT Codex Connector

   Last edited by novkien

   OwnerAuthor

   More actions

   Nghiên cứu bổ sung V100/SM70: kỹ thuật có thể chuyển vào llama.cpp-fork, phần đã có và kế hoạch kiểm chứng

   Ngày rà soát: 2026-10-01. Đã đọc issue, kiểm tra lại master = e58b770f3c2346deab036d9e4c3b2ea61ca0bbc5, đối chiếu các đoạn CUDA/graph liên quan với tài liệu thiết kế, mã nguồn được dẫn và PR của các dự án tối ưu V100.

   Kết luận: còn nhiều cơ hội ngoài năm nhóm trong body, nhưng phần đáng chuyển là dataflow, layout và lựa chọn kernel theo shape, không phải thay engine hoặc bê nguyên kernel FP8/NVFP4 sang Q5_K. Những bổ sung có giá trị nhất là: backend-only weight repacking; GEMM tuning tách khỏi layout; fused/tiled dequantization cho prefill; tối ưu việc đọc/chuyển đổi KV; phân vùng attention theo KV đang sống; ổn định CUDA Graph; và tối ưu đúng đường GDN đang được dispatch.

   Giới hạn bằng chứng: đây là nghiên cứu nguồn và thiết kế tích hợp. Không có benchmark GPU mới trên máy chủ trong lượt nghiên cứu này; không sửa runtime, đổi cấu hình, deploy hay xác nhận mức tăng tok/s. Những con số từ donor bên dưới là tác giả donor báo cáo, không phải kết quả của fork này.

   1. Baseline và những điều không được suy diễn sai

   Giữ nguyên baseline của issue: 4x V100-SXM2-16GB, GGUF UD-Q5_K_XL, K/V Q8_0, architecture qwen35, D256/GQA6, b=2048, ub=1024, 4 slots, layer split hiện tại và MTP tối đa 2 draft. Mỗi slot có thể cần xác minh tối đa 3 token; tổng số row của GEMM phải lấy từ trace thực tế, không đồng nhất số slot với query length của attention.

   Đọc lại code cho thấy:

   Thành phần	Fork đã có	Delta cần tìm, không làm lại từ đầu
   CUDA Graph	V100 được phép dùng graph; chỉ GPU có CC thấp hơn Volta bị loại tại gate kiến trúc. Khi properties thay đổi, warmup bị reset. F1	Tần suất cache hit, reset, recapture và nguyên nhân shape/address thay đổi.
   Volta cuBLAS	Với operand FP16, nhánh Volta chủ động chọn compute/output FP32. F2	Giảm conversion/scratch hoặc chọn kernel phù hợp hơn, không hạ toàn bộ accumulation xuống FP16.
   GDN recurrent	Kernel giữ state shard và Q/K trong register, lặp qua token, hỗ trợ snapshot phục vụ rollback. F3	Không đề xuất “đưa state vào register” như tính năng mới. Đo register pressure, input preparation, snapshot và dispatch.
   GDN prefill	delta-net-base.cpp đã có cả fused và chunked graph; chọn theo fused_gdn_ar/ch. Chunking GDN thường dùng CS=64; KDA có nhánh riêng. F4	Không kết luận từ TODO trong CUDA kernel rằng toàn bộ engine chưa có chunked GDN.
   QKV execution	Backend đã có logic nhận diện fan-out attn_norm với ba nhánh, tổ chức concurrency và bảo vệ lifetime. F5	Shared input staging/fusion chỉ là bổ sung nếu đường hiện tại còn lặp conversion hoặc launch; thêm stream không tự động tốt hơn.
   GQA	Nhánh Volta chọn nhóm 8/4/2/1 theo tính chia hết. Với GQA6, đường đủ điều kiện hiện chọn nhóm 2, không phải nhóm 6. F6	Khai thác reuse giữa đủ sáu query heads, với contract riêng cho Q8_0.
   Sparse attention	Đã có mask compaction, nhưng fast path MMA sparse đang có điều kiện turing_mma_available. F6	Không cho rằng hạ gate là đủ để chạy SM70; không biến full attention thành approximate sparse attention.

   Qwen35 còn chứa các loại layer khác nhau; tối ưu D256 full attention không đồng nghĩa tối ưu toàn bộ mạng. Cần cộng thời gian từng nhóm layer/operator trước khi chọn việc. F7

   2. Donor nào đáng học và học phần nào?

   Donor	Bằng chứng đáng đọc	Cách dùng cho fork
   zhinianqin/marlin_v100	Workspace Marlin dense/MoE dành SM70; policy CTA/warp shape, split-K, metadata cache, tách PackedMacroN khỏi CTA_N. Snapshot kiểm tra: 3f16d442cdb4c24dd225bbec196c982a54d9a31c. S1, S2	Học thiết kế iterator/layout và tuning. Không dùng decoder AWQ/FP8 thay decoder GGUF.
   dnv2003/v100-skinny	QPN fragment ordering, weight reuse giữa nhiều row, giới hạn khác nhau giữa NVFP4 và FP8; ghi rõ overflow, fallback và chi phí các bản repack. S3, S4	Học cách đưa dữ liệu nén vào operand Volta ở điểm tiêu thụ; bảo toàn mã và scale GGUF.
   geoffwatts/ninfer-v100	Repack tại load có runtime layout tag; reuse activation staging giữa projection; chia verification width thành các nhóm đã được qualification. Snapshot kiểm tra: b37d0dd3e1163b9d802d8bccfa89918bf68d793e. S5	Học ownership, lifetime, shape specialization. Không sao chép format KV riêng của NInfer.
   1CatAI/1Cat-vLLM	PR ggml-org#268: page-ID reuse và aligned wide loads; ggml-org#328: bỏ các tile chắc chắn nằm ngoài sliding window có sẵn; issue ggml-org#441: tile kiến trúc mới có thể rất tệ trên pre-Ampere. S6, S7, S8	Học cách giảm instruction/address/conversion overhead, chọn tile riêng SM70 và chứng minh route thực sự chạy.
   mistrjirka/llama.cpp	Các donor native GGUF/Q8_0 đã nêu trong body: MMVQ Q5_K/Q6_K và attention Q8 Volta. S9	Ưu tiên làm control/candidate gần fork nhất trước khi thiết kế layout mới. Vẫn phải giải quyết W=3 và multi-slot thay vì chỉ W=4.
   ai-bond/flash-attention-v100	FlashAttention-2 viết cho V100/Titan V. S10	Nguồn tham chiếu bổ sung về attention SM70, không phải thư viện có thể gắn thẳng vào tensor/KV layout ggml.
   NVIDIA CUTLASS và FasterTransformer	Double-buffered GEMM; fusion; workspace reuse; tìm cấu hình GEMM theo shape và lưu kết quả. S11, S12	Dùng nguyên lý có nguồn chính thức để thiết kế CUDA native, không thêm dependency PyTorch/vLLM vào runtime.

   Không tính các fork NInfer có cùng phần kế thừa là nhiều bằng chứng benchmark độc lập. NInfer cũng ghi rõ một đường long causal prefill dùng kernel FlashAttention MMA lấy từ llama.cpp: phần này có thể là vòng quay mã nguồn, không phải phát minh mới để “port ngược”. S5

   3. Các hướng bổ sung chi tiết

   A. Offline shape tuning và dispatch riêng SM70 — ưu tiên P0/P1

   Kỹ thuật donor: Marlin V100 mô tả geometry gồm CTA_M, CTA_N, CTA_K, Warps, WarpM, WarpN, WarpK, cùng split-K và lựa chọn cách cache metadata. Có cả kiểm tra phân rã warp lẫn contract iterator/epilogue; tổ hợp đúng số học vẫn có thể không hợp lệ với layout. FasterTransformer dùng bước tìm GEMM config trước khi chạy. S2, S12

   Thiết kế đề xuất:

   1. Thu một shape census từ workload hiện có: tensor name/type, ne[], nb[], GPU, GEMM M/N/K, số sequence, query width, live KV, kernel/nhánh được chọn, conversion bytes, scratch peak. UD-Q5_K_XL là tên gói model, không đủ để kết luận mọi tensor là Q5_K.
   2. Tách control theo các vùng: decode 1 row; verification thực tế 2/3 row; các số row phát sinh từ 2/4 slots; prefill 1024 và các tail. Xác định M/N/K theo quy ước toán học thống nhất vì layout ggml có thể đặt trục khác tên donor.
   3. Chạy offline một tập nhỏ cấu hình hợp lệ: MMVQ hiện tại, donor native, cuBLAS hiện tại, candidate mới; chỉ thêm MMQ/SM70 Tensor-Core path khi đã xác định nó xử lý đúng quant type đó.
   4. Lưu dispatch table theo capability + quant type + shape/alignment + layout version. Dùng ngưỡng đơn giản sau khi đo, không copy các mốc M=8/16/96 của donor.
   5. Runtime chỉ lookup; không benchmark hàng loạt trong request đầu tiên. Có fallback hiện tại và switch debug để A/B. Không bắt mọi V100 phải dùng cùng tile nếu register/shared-memory profile khác.

   Vị trí: ggml-cuda.cu matmul dispatch, mmvq.cu, mmq.cu/helpers liên quan, test/benchmark backend hiện có.

   Điều kiện thành công: bảng tuning phải giúp các shape có trọng số thời gian đáng kể, không chỉ chọn shape đẹp. Bao gồm chi phí conversion, reduction và workspace trong thời gian op. Nếu kernel chỉ nhanh khi loại bỏ conversion khỏi phép đo thì chưa đạt.

   B. Lossless backend-only weight repacking — ưu tiên P2, có thể có lợi lớn nhưng không miễn phí

   Kỹ thuật donor: QPN sắp mã weight theo register fragment order để bớt shuffle trong inner loop. NInfer biến đổi code/scale planes khi load, gắn tag cho view runtime, rồi giải phóng temporary. Marlin tách macro-layout của packed weight khỏi tile GEMM đang chạy. S3, S5, S2

   Áp dụng cho GGUF: không requantize, không sửa GGUF trên đĩa; chỉ đổi thứ tự biểu diễn backend của những tensor đủ điều kiện. Decoder phải xử lý đúng từng loại Q5_K/Q6_K và loại thực tế trong census, bao gồm scale/min/high bits, không coi chúng là một format INT4 đơn giản.

   Contract cần thiết trước khi code:

   • Layout tag/version thuộc backend buffer hoặc metadata có ownership rõ; không thay bytes dưới một tensor vẫn bị consumer hiểu là layout GGUF gốc.
   • Repack phải đảo ngược được và kiểm tra byte parity với mã/scale nguồn. Byte parity của repack không đồng nghĩa kết quả phép nhân mới tự động bitwise bằng kernel cũ.
   • Quyết định rõ một layout dùng chung cho PP/TG hay hai bản packed. v100-skinny mô tả trường hợp NVFP4 giữ cả bản QPN lẫn Marlin: không được gọi mọi prepack là “zero-extra-VRAM”. S3
   • Unsupported shape, CPU access, tensor copy, view, readback, cuBLAS fallback đều phải hiểu layout hoặc có bước unpack có giới hạn. Không giải phóng bản gốc rồi để fallback đọc nhầm.
   • Shared/tied weights giữa target và MTP chỉ repack một lần khi cùng representation/device; không double-free hoặc tạo bản thừa âm thầm.
   • Các adapter/LoRA và tensor có mutation cần rule invalidation hoặc loại khỏi eligibility ban đầu.
   • Repack ngoài graph capture. Đo cả peak khi cùng tồn tại source + destination + temporary; cân nhắc xử lý theo từng tensor trên 16GB.

   Vị trí: CUDA buffer/tensor upload trong ggml-cuda.cu, metadata ở CUDA backend, kernel đọc quantized weight; chỉ đụng loader chung nếu backend hiện tại không diễn đạt đủ ownership. Ưu tiên không sửa định dạng GGUF hay API người dùng.

   Tiêu chí bỏ hướng: thêm một bản nén làm workload target/draft hết VRAM, hoặc fallback phải unpack quá thường xuyên đến mức không còn lợi ích E2E.

   C. Prefill: fused/tiled GGUF dequantization + GEMM thay vì dựng toàn bộ tensor FP16 — ưu tiên P2

   Trong ggml_cuda_mul_mat_cublas_impl, khi type không trùng compute type, có allocation/conversion cho các phần tử của tensor đầu vào của phép gọi trước khi GEMM. Đây là một nơi có thể phát sinh conversion/scratch; chưa có trace để kết luận chiếm bao nhiêu PP. F2

   Đề xuất: giữ weight nén trong HBM, đọc một tile, giải mã sang FP16 ở register/shared memory gần consumer, dùng Volta MMA với accumulator FP32. So sánh hai thiết kế: đọc GGUF gốc trực tiếp, và đọc layout backend của mục B. Không bắt buộc phải repack mới có prototype.

   Tổ chức software pipeline hợp SM70: prefetch tile kế tiếp vào register khi xử lý tile hiện tại, double-buffer shared memory khi thực sự giảm stall, bố trí code/scale để coalesced load và tránh bank conflict. CUTLASS mma_pipelined.h là mẫu double-buffered threadblock GEMM để học. S11

   Các lưu ý quan trọng:

   • Không bê pipeline phụ thuộc tính năng kiến trúc GPU mới vào V100 chỉ bằng cách đổi compiler arch.
   • Double buffer tăng live registers/shared memory; giảm stall nhưng làm spill hoặc mất CTA residency có thể chậm hơn.
   • Scale cache cũng dùng register; phải đo vector_words so với chia theo lane thay vì giả định giữ mọi scale là tối ưu.
   • Cần threshold PP/TG riêng. GEMM lớn của cuBLAS có thể thắng custom kernel; giữ đường cuBLAS cho các vùng thắng.
   • Không dùng FP16 accumulation toàn K chỉ để tăng điểm microbench. Control Volta hiện chọn FP32 có chủ đích. F2

   Bộ nhớ: nếu một phép gọi đang tạo tensor FP16 KxN thì riêng phần đó có kích thước 2*K*N byte; đây là công thức cho tensor của phép gọi, không phải khẳng định engine giải nén toàn model cùng lúc. Candidate phải báo cả dequant tile, split-K partials và output staging, không chỉ bỏ một khoản khỏi báo cáo.

   D. Attention: tối ưu đường đọc KV và dependency trước khi thay toàn bộ thuật toán — ưu tiên P1/P2

   PR 1Cat ggml-org#268 là bằng chứng cụ thể: reuse page IDs và ghép các lần load/chuyển đổi đã cải thiện batched E5M2 XQA. Tác giả báo operator khoảng 1.21–1.27x ở các shape mẫu, trong khi E2E pure decode ở hai cell long-context tăng khoảng 6–8%. Counter cho thấy HBM bytes và occupancy gần như không đổi; giảm instruction/load-request/dependency mới là cơ chế chính. S6

   Với Q8_0 của fork:

   1. Đo riêng address calculation, đọc scale/code, chuyển sang FP16, publish shared, QK, softmax, PV, reduction. HBM chưa bão hòa không có nghĩa thiếu bandwidth phần cứng; consumer có thể đang chờ conversion hoặc dependency.
   2. Hoist phép tính địa chỉ/stride/scale ra ngoài inner loop khi hợp lệ. ggml không nhất thiết có paged-KV page table như vLLM: chuyển nguyên lý reuse metadata, không thêm page table chỉ để giống donor.
   3. Vectorize code loads khi alignment thực sự cho phép. Quant block chứa cả metadata và payload, nên không được mặc định mọi đầu block hoặc mọi row đều aligned 16 byte. Kiểm tra type layout, offsets và nb[], có scalar/tail path.
   4. Reuse một tile K/V cho nhiều query head chung KV và nhiều verification query; đo riêng W=1,2,3. Có thể giữ query ở tile nội bộ lớn hơn, nhưng row không hoạt động phải được mask hoàn toàn và không phát sinh token/state thật.
   5. Xử lý cả long-prefix suffix PP và cold PP. Các shape này có tỷ lệ query/KV và masking khác nhau; một nhánh không bắt buộc tối ưu cả hai.

   Đây là phần mở rộng thực chất cho D256/GQA6 trong body. Không đổi Q8_0 sang E5M2/FP16 để làm donor kernel chạy được; làm vậy phá baseline và thay bài toán bộ nhớ.

   Vị trí: fattn.cu, fattn-common.cuh, các MMA/vec/tile helpers và specialization mới có gate hẹp. Dùng FLASH_ATTN_EXT hiện có.

   E. Split-KV theo live work, bounded scratch và bỏ chính xác tile bị mask — ưu tiên P1 nếu trace chứng minh khoảng trống

   v100-skinny mô tả việc loại chi phí decode phụ thuộc quá mạnh vào max_model_len khai báo khi KV thực tế còn ngắn. 1Cat ggml-org#328 nhảy thẳng tới tile đầu nằm trong sliding window vốn có của DFlash2 draft, thay vì đi qua prefix chắc chắn bị mask. S3, S7

   Cần kiểm tra trong fork, chưa coi là bug đã xác nhận:

   • Split count/launch grid/partial reduction đang phụ thuộc allocation capacity, KV length của op hay số token thật được mask cho từng sequence?
   • Có scan/clear/reduce vùng scratch vượt quá số partition thực sự có dữ liệu không?
   • Khi nhiều sequence có độ dài lệch nhau, có vô tình tính mọi sequence như sequence dài nhất không?

   Thiết kế: chọn số partition từ query width, số KV head, live/visible KV range, SM budget và workspace cap. Ít partition quá sẽ thiếu parallelism; nhiều quá tăng launch/partial/reduction. Không hardcode “luôn split tối đa”.

   Để giữ graph ổn định, có thể dùng một số bucket partition và bounds trên device; chỉ áp dụng nếu phần work thừa trong bucket nhỏ hơn chi phí recapture. Không lấy chiều dài về CPU bằng synchronous readback mỗi token chỉ để tiết kiệm một kernel.

   Causal prefix hoặc sliding window phải dựa trên mask/position thật. Full attention của target vẫn phải thấy toàn bộ KV hợp lệ; tuyệt đối không đưa cửa sổ 2048 của DFlash2 vào model hiện tại. Với mask tùy ý phải chứng minh toàn tile bị loại, không chỉ nhìn query cuối.

   FP32 partial online-softmax merge phải xử lý m/l/output nhất quán; một partition rỗng hoặc hàng all-masked không được sinh NaN hay đọc dữ liệu cũ. Bộ nhớ partial tăng theo số query, partition và head dimension; cấp phát có giới hạn theo GPU hiện có.

   F. CUDA Graph ổn định theo shape, không chỉ “bật graph” — ưu tiên P1 nếu host/launch gap đáng kể

   Fork đã có graph key/cache và yêu cầu các lần gọi có properties ổn định để qua warmup; properties đổi sẽ quay lại execute trực tiếp. F1 NInfer là donor hữu ích về tổ chức shape-specific execution và dữ liệu runtime, không phải bằng chứng rằng fork chưa có graph. S5

   Delta đề xuất:

   • Đếm hit/replay, first capture, update, reset và lý do reset. Tách target decode, target verify, draft, PP; không dùng một phần trăm graph-hit chung che một pha luôn miss.
   • Kiểm tra key/properties có trộn dữ liệu thay đổi mỗi token với topology/shape hay không. Chỉ bỏ một trường khỏi key khi dữ liệu đó đã được cập nhật bằng cơ chế graph-safe tương đương.
   • Cố định ownership/address của metadata và scratch; cập nhật positions, lengths, slot map trên device hoặc qua async copy đúng stream dependency.
   • Nếu cần thêm shape buckets, bắt đầu từ shape quan sát được, có eviction/lifetime rõ. Không capture mọi tổ hợp context x slots x draft; graph instances và pinned buffers cũng tốn bộ nhớ.
   • Không gộp query của các sequence khác nhau theo cách làm sai mask. Không padding batch logic để vô tình cập nhật GDN/KV của slot không hoạt động.
   • Tránh allocation, host sync, readback và thay đổi buffer lifetime trong capture; giữ fallback an toàn khi capture không hợp lệ.

   Vị trí: graph key/properties/cache và compute/capture trong ggml-cuda.cu, scheduler/input plumbing liên quan nếu trace yêu cầu. Đây là công việc nhỏ hơn GPU-resident MTP toàn tuyến, nên có thể làm trước nếu nó tự mang lại lợi ích.

   G. Shared activation staging và fusion đúng chỗ — ưu tiên P1/P2

   NInfer reuse cùng staged FP16 activation cho nhóm projection. FasterTransformer gộp các operator nhỏ và tái sử dụng workspace giữa layer. S5, S12 Fork đã có concurrency/fusion infrastructure, nên chỉ thêm phần thiếu được trace xác nhận. F5

   Các candidate cần đo:

   • Q/K/V hoặc gate/up cùng dùng input: có đang convert/quantize input lặp không? Reuse một bản input staging khi dtype/layout/numerics giống nhau; nếu khác phải giữ buffer riêng.
   • Fuse epilogue hợp lệ như bias/activation/gate/residual khi tránh được intermediate HBM traffic và không phá consumer khác. RMSNorm của Qwen không được thay bằng LayerNorm của donor.
   • Với Qwen35 GDN có projection ghép sẵn wqkv; không “merge QKV” thêm lần nữa như một tối ưu mới. F7
   • Đánh giá fusion so với QKV multi-stream hiện tại: một kernel lớn hơn có thể giảm launch nhưng tăng register pressure và làm mất overlap. Chỉ giữ phương án thắng E2E.
   • Workspace có thể dùng lại giữa các op/layer không đồng thời; target, draft và các stream đang chạy cùng lúc không được alias chỉ vì có cùng kích thước.

   Vị trí: CUDA graph optimizer/fusion matcher, norm.cu, matmul epilogues và graph builder tương ứng, sau khi xác định pattern thật. Cập nhật allocation dependencies cùng lúc với fusion, không tách thành bước “sửa sau”.

   H. GDN: tối ưu đường prefill thực sự, không chỉ kernel recurrent — ưu tiên P1 để profile, P2 để implement

   Đây là phần cần bổ sung mạnh vào kế hoạch vì code có hai đường khác nhau:

   • build_delta_net_fused() tạo GGML_OP_GATED_DELTA_NET; kernel CUDA hiện giữ state trong register và đi tuần tự theo token. F3, F4
   • build_delta_net_chunking() tạo các graph op cho pad, transpose/contiguous, cumsum, decay mask, GEMM và triangular solve; lựa chọn phụ thuộc fused_gdn_ch. F4

   Thực hiện nghiên cứu đúng thứ tự:

   1. Ghi fused_gdn_ar/ch, số token/sequence và kernel/operator trace cho actual PP và MTP verify. Các flag hiện tại là baseline, không đổi chúng trong A/B của một thay đổi khác.
   2. Nếu fused sequential path chiếm phần lớn PP, thiết kế native chunked SM70 candidate hoặc crossover policy có bằng chứng; không thay decode/3-token verify bằng một pipeline chunk lớn có overhead cao.
   3. Nếu chunked graph đang chạy, tìm các materialization có thể bỏ: transpose -> cont -> cumsum, lặp decay mask, scale/beta preprocessing và nhiều kernel nhỏ quanh GEMM/solve. Có thể prototype fused prefix-scan/preprocessing với output layout phù hợp consumer để tránh copy trung gian.
   4. Cân nhắc GEMM/triangular-solve theo tile SM70 và reuse scratch. Không giả định mọi tensor đầu vào FP32 có thể đổi sang FP16 an toàn để bật Tensor Core; cần qualification riêng cho phép biến đổi precision đó.
   5. Giữ đường register-resident hiện tại cho short query nếu nó thắng. Snapshot chỉ được bỏ khi chứng minh không cần cho acceptance/rollback; không lấy giảm số lần ghi checkpoint làm “speedup” bằng cách đổi semantics.

   Vị trí: src/models/delta-net-base.cpp, ggml/src/ggml-cuda/gated_delta_net.cu, CUDA cumsum/solve_tri/matmul helpers, state layout hiện hành.

   Đây là đề xuất tích hợp suy ra từ code fork và nguyên lý fusion/tiling, không phải tuyên bố đã tìm được một kernel GDN donor tương thích hoàn toàn. Không cần đưa Torch/TileLang vào ggml chỉ vì donor dùng chúng.

   I. MTP: tối ưu hidden-state handoff và phần control/sampling theo contract — ưu tiên P2, sau khi có attribution

   Giữ hướng device-resident MTP trong body, nhưng chia thành các khoản đo riêng: lấy hidden state, copy/compact row, target/draft synchronization, draft LM head, sampling, verification và rollback. Nhanh một kernel draft mà mất nhiều hơn ở synchronization không phải thắng.

   Candidate nhỏ nhất là handoff buffer device-to-device/device-view có event và lifetime rõ, giữ public host getter cho caller cần CPU. Không âm thầm đổi API từ trả host pointer sang device pointer. Handoff cross-GPU phải tuân thủ P2P policy; same-device và peer/non-peer là ba trường hợp khác nhau.

   Có thể học device-side selection/sampling của các native engine và FasterTransformer, nhưng không khẳng định fork thiếu toàn bộ GPU sampling chỉ vì thấy CPU code. Audit sampler chain thực tế trước. Grammar, penalties, logprobs, RNG, cancellation và MTP accept/reject vẫn phải đúng; GPU argmax chỉ tương đương khi semantics đang thực sự là greedy. S12

   NInfer còn có context-lookup proposals được target xác minh. Đây là hướng riêng cho workload lặp/copy, không nằm trong baseline MTP=2: không kéo draft depth lên 15 để lấy kết quả của donor. Tương tự, tối ưu/thu gọn proposal LM head chỉ có thể áp dụng phía draft khi đã định nghĩa rõ proposal distribution; tuyệt đối không cắt vocabulary của target để tăng tốc. S5

   J. Communication/topology: chỉ bổ sung overlap có dependency thật — ưu tiên P2/P3

   Không lặp lại đề xuất “bật NVLink/all-reduce” vì fork đã có. Việc đáng nghiên cứu là timeline compute/copy/event tại ranh giới layer và giữa target/draft; dùng event/stream dependency thay cho host-wide sync khi tương đương, pooled pinned staging cho đường không-peer, chunked transfer nếu giúp overlap mà không tăng launch quá nhiều.

   Giữ nguyên allowlist #3/PR #4. Không enable P2P cho các cặp PCIe chưa an toàn, không suy từ GPU0 sang mọi GPU và không nhầm CUDA ordinal theo UUID với physical ordinal. Thêm stream không tạo parallelism cho các layer vẫn có dependency tuyến tính. Muốn tensor parallel + Q8 KV cần thiết kế meta-device/cache-layout riêng như body đã cảnh báo, không gộp vào PR kernel nhỏ.

   4. Những hướng phải loại khỏi baseline hoặc ghi rõ giới hạn

   • Không lấy mã FP8/NVFP4 làm bằng chứng V100 có native FP8/NVFP4 Tensor Core. Donor giải mã sang operand mà SM70 hỗ trợ; representation trên HBM và kiểu phép nhân là hai chuyện khác nhau. S3, S5
   • Không copy preset GPU mới. Issue 1Cat Eliminate ggml_forward_mul_mat_xxx() branch for non-contiguous src0 ggml-org/llama.cpp#441 báo pathology từ tile tuned trên GB300 cho pre-Ampere và còn chỉ ra đã có lúc tune nhầm file không được dispatch. Model/operator của họ khác workload này: học cách kiểm tra route, không mang hệ số 19x vào dự báo cho Qwen35 của fork. S8
   • Không đổi Q8_0 thành INT8 group-64 + Hadamard của NInfer. Đó là quant/layout/transform contract khác, cần issue/approval và đánh giá chất lượng riêng. S5
   • Không áp dụng sliding window/sparse pruning gần đúng cho full attention. Chỉ bỏ work đã bị mask theo semantics gốc. S7
   • Không mặc định FORCE_MMQ, FORCE_CUBLAS hay FP16 accumulation là luôn tốt. Mỗi đường có trade-off shape, numerics và VRAM; đây là đối tượng đo, không phải kết luận tối ưu.
   • Không tính một PR có nhãn Merged là một optimization độc lập. Ví dụ Add requirements to readme ggml-org/llama.cpp#269 có lịch sử audit zero-diff rồi merge để ghi nhận main đã chứa phần hiệu lực; không cộng speedup Add requirements to readme ggml-org/llama.cpp#269 vào Failing CI autobuilds ggml-org/llama.cpp#268 hay lấy title thay cho diff. S13
   • Không gom mọi thay đổi vào một giant PR. Prefix/cache retention liên quan [cache] LCP slot reuse discards another conversation without RAM preservation + lookup predicate gap (tracks novkien/llama-proxy#416) #1, selective P2P liên quan Research: selective CUDA peer access for four V100s with unsafe PCIe P2P #3, phase exclusion liên quan Research: native cross-slot prefill/TG phase exclusion #5; giữ ranh giới ownership. Vision encoding là workload riêng nếu mở rộng về sau, không dùng điểm PP text để chứng minh cải thiện vision.

   5. Kế hoạch thực nghiệm và điều kiện nghiệm thu bổ sung

   5.1. Evidence pack trước khi chọn kernel

   Lưu commit/binary hash, driver/toolkit/compiler, build flags, model và quant metadata, actual tensor types, GPU UUID/ordinal mapping, topology, memory headroom, graph config, GDN route và MTP depth. Không suy deployed revision chỉ từ physical checkout HEAD đã được issue cảnh báo.

   Thu trace trên đúng cấu hình baseline, với profiling tắt khi lấy số E2E cuối. Các counter/trace nên có chế độ opt-in; tránh log đồng bộ theo từng token làm sai latency. Kết quả cần chỉ ra tên candidate, selected call count, fallback call count/reason và tỷ trọng thời gian của đường đủ điều kiện.

   5.2. Ma trận test tối thiểu

   Nhóm	Các trường hợp phải có
   Matmul	Từng quant type thực tế; decode M=1; verify M=2/3; shape thật của 2/4 slots; PP ub=1024 và tail; alignment/stride không thuận; output head riêng.
   Attention	D256/GQA6/Q8_0; W=1/2/3; one/multi-sequence; ragged live lengths; cold PP và cached-prefix suffix; query/KV tail, mask lệch vị trí.
   GDN	Fused và chunked routes được xác định; one token, three-token verify, PP, chunk boundary/tail; output lẫn final/intermediate state. Không thay cả route lẫn kernel trong cùng A/B mà không tách attribution.
   MTP/state	Accept 0/partial/all, reject sớm/muộn, rollback rồi tiếp tục decode, reuse slot, cancel, EOS, checkpoint restore, các sequence độc lập.
   Graph/lifetime	Replay lặp lại; đổi slot count/query width; graph miss/fallback; reuse/free buffer; concurrent target/draft/stream; unload/reload không leak.
   Memory	Peak trên từng GPU, không chỉ tổng 64GB; weights + KV + GDN snapshots + target/draft + graph + all scratch, bao gồm transient repack/conversion.
   Fallback	Non-SM70 không bị đổi behavior; unsupported type/head dimension/stride/mask chạy đường cũ; candidate tắt trả về control thật.

   Khung timing giữ lại fixture 8192/32768 và slot-order của #5 để đối chiếu được, bổ sung decode/MTP và cached prefix. Các context dài hơn chỉ chạy trong ngân sách VRAM đã kiểm tra; không có yêu cầu cố chạy một case sẽ OOM. Trộn prompt ít lặp, code, prose và nhiệm vụ copy để tránh speculative acceptance bị thổi phồng bởi filler.

   5.3. Numerical, memory và synchronization gates

   • Repack round-trip exact ở mức code/scale; output kernel so với native/reference với absolute + relative error đã định trước theo op, và dữ liệu activation thật có outlier. Không đặt tolerance sau khi nhìn kết quả candidate.
   • Attention kiểm tra causal/sequence mask, GQA head mapping, all-masked/empty partition, softmax stability; tile/reduction đổi thứ tự có thể không bitwise nhưng phải qua tiêu chí đã định.
   • GDN kiểm tra cả trạng thái để tiếp tục sinh token sau rollback, không chỉ output của một forward.
   • Sampling/quality: giữ sampling config thật; deterministic probes khi áp dụng được, kèm task/quality tests và logprob/distribution checks phù hợp. Một đoạn trả lời nghe hợp lý không phải proof.
   • Theo NVIDIA, Volta có independent thread scheduling: dùng warp synchronization/mask và barrier đúng contract, không dựa vào lockstep ngầm. Chạy memcheck/racecheck/synccheck phù hợp cho candidate. S14
   • Shared-memory/register/stack/spill phải được kiểm tra từ build/profiler. SM70 có ngân sách shared memory 96 KiB/SM; allocation động vượt ngưỡng 48 KiB cần opt-in phù hợp. Nhiều shared memory hơn không tự động nhanh hơn. S14

   5.4. Cách kết luận có lợi

   Dùng warmup, các lượt A/B xen kẽ hoặc đảo thứ tự, ít nhất 5 cặp đo cho các cell quyết định; tăng số lần nếu nhiễu lớn. Báo median, độ phân tán và khoảng tin cậy phù hợp. Lưu cả case thua.

   Báo operator time có đủ conversion/reduction và E2E riêng: cold PP, prefix-suffix PP, TTFT, TG/TPOT, committed tokens/s, MTP acceptance và peak VRAM. Không đếm draft/rejected tokens là output throughput. GPU utilization cao hơn không phải tiêu chí merge.

   Áp dụng Amdahl để tránh hứa quá mức: nếu một op chiếm tỷ lệ f, nhanh lên s lần thì trần tăng tốc tổng lý tưởng là 1 / ((1-f) + f/s), chưa tính overhead mới. Ví dụ giả định f=20%, s=2 thì tổng chỉ khoảng 1.11x, không phải 2x. Tỷ lệ f của máy này hiện chưa đo trong lượt nghiên cứu này.

   6. Thứ tự triển khai đề xuất cho một scope được phê duyệt

   1. P0: route/shape/time/memory census. Kết quả là một báo cáo có attribution, không thay behavior.
   2. P1: một thay đổi hẹp gần code hiện tại. Chọn theo census: native MMVQ W=3; Q8 attention wide-load/reuse; live partition; hoặc graph miss. Không làm cả bốn để rồi không biết cái nào giúp.
   3. P2: layout hoặc pipeline mới. Chỉ khi P1 cho thấy dequant/GEMM hay GDN là giới hạn: repack có ownership/fallback, fused prefill, hoặc GDN chunk/fusion. Tách từng candidate để rollback được.
   4. P2/P3: MTP dataflow và transport. Chỉ xử lý sau khi timeline chứng minh host/copy/communication còn là critical path; tensor mode + quantized KV là thiết kế lớn riêng.

   Checklist cho PR được chọn:

   • Source SHA và diff donor cụ thể; chỉ rõ phần kế thừa, phần tự thiết kế; kiểm tra license/NOTICE theo từng file. v100-skinny phân biệt MIT kernel với Apache-2.0 fork patches; không áp một license chung tùy tiện. S3
   • Contract eligibility + fallback + layout/ownership + precision + workspace cap được viết trước.
   • Prove actual dispatch trên workload hiện tại, không dùng thay đổi model/KV/MTP depth để làm kernel đủ điều kiện.
   • Backend correctness và real-activation/outlier tests đạt; unsupported route không regress.
   • Graph, state, rollback, concurrency và lifetime tests đạt.
   • Không OOM/leak trên từng V100 16GB, cả peak load lẫn steady inference.
   • E2E benefit lớn hơn nhiễu; báo cả neutral/regression, không chỉ peak microbenchmark.
   • Có kill switch/fallback, review source và tái kiểm tra với revision mới trước merge.

   Đề xuất quyết định: ưu tiên đo và thử các delta native trước; coi repack + SM70 GEMM là khoản đầu tư lớn hơn, không phải patch vài dòng. Issue nên tiếp tục mở: comment nghiên cứu không thay cho implementation hay xác nhận tối ưu thành công.

   Các liên kết F1–F7 trong nội dung trỏ tới code fork được pin theo commit; S1–S14 trỏ tới nguồn tham khảo tương ứng. Với nguồn theo branch hoặc repo homepage, phải pin lại revision và kiểm tra diff trước khi port code.
2. novkien commented on Oct 1, 2026

   

   novkien

   on Oct 1, 2026 – with ChatGPT Codex Connector

   OwnerAuthor

   More actions

   Simulation results: PP/TG shortlist, numerical tests, and rejected shortcuts

   Research date: 2026-10-01. Owner direction: continue the V100 research and actually test candidates in a self-created simulation environment. Rechecked master before publication: e58b770f3c2346deab036d9e4c3b2ea61ca0bbc5.

   Conclusion: the strongest first native experiment is direct-Q8_0 attention for SM70 / D256 / GQA6 / W3 verification, avoiding full-view FP16 materialization and reusing KV across query/head rows. Long-prefix PP needs its own larger query tile. The simulation rejects both “force the existing vector route and assume it wins” and “reuse a tiny decode tile for all prefill.”

   This is a workload-specific prioritization, not a measured V100 speedup. Overall target-performance verdict: PARTIAL.

   1. What actually ran

   The sandbox has 5 logical CPUs, Python 3.13.5, NumPy 2.3.5, CPU-only PyTorch, no nvcc, no nvidia-smi, and zero CUDA devices. No target model weights, CUDA binary or GPU instruction trace was loaded. No production source/configuration/service/GPU job was changed.

   Executed an independently authored CPU semantic harness and a deterministic source-informed cost/resource simulator:

   • 273 numerical checks passed, including 2 expected negative controls.
   • 17 analytical self-checks passed.
   • 2,400 attention sensitivity scenarios, 1,680 GEMM sensitivity scenarios, and 72 source-tiling/cache scenarios were evaluated. These are counterfactual parameter combinations, not GPU runs or statistical evidence of success probability.
   • Re-ran the complete reproduction script; all 11 result JSON files reproduced exactly after excluding CPU wall-clock diagnostics.

   Numerical suite	Checks	Coverage	Maximum absolute error against its independent reference
   Q5_K/Q6_K core	52	Scalar/vector block decode, reversible byte-plane packing, tiled FP16-operand GEMM	1.013e-4
   Q8_0 attention core	111	D256/GQA6, W1/2/3/4/8/16, tails, causal/ragged-isolated masks, split reduction; two negative controls	1.858e-6 for correct variants
   GDN core	102	Chunk algebra versus recurrent reference, final state, accepted-prefix state	4.977e-7
   Long-span attention	1	W3, N131072, D256, Hq24/Hkv4, tile32, 32 partitions	1.047e-8
   PP-sized attention	1	W1024, N2048, D256, Hq24/Hkv4, online/split algebra	1.753e-6
   K5120 GEMM	4	Q5_K/Q6_K, M3 and M1024, 64 output rows	1.647e-4
   T1024 GDN	2	D128, chunk64, outputs and final state	2.014e-7

   Attention/GDN tolerances were atol=2e-5, rtol=2e-5; GEMM atol=2e-4, rtol=2e-5. The predicate is max_abs_error <= atol + rtol * max_abs_reference, with finite outputs. Block decoding and packing round-trip require exact equality.

   The negative controls deliberately omit a sequence mask and overflow an unscaled FP16 partial. Both are rejected. The overflow is not a reproduction of a donor bug: current kernels already contain scaling safeguards.

   Limit: these CPU prototypes do not validate native HMMA fragments, per-instruction rounding, MMVQ Q8_1 activation quantization, existing FP16 probability/FTZ intermediates, CUDA races, or complete-model quality. The repack test is an invertible byte-plane demonstration, not a tuned QPN layout. W1024 CPU algebra is not a proposal to allocate all queries in one CUDA block. CPU elapsed times are not extrapolated into V100 tok/s.

   2. New concrete source finding: W3 crosses a KV-materialization boundary

   In the current dispatcher, GQA6 yields effective power-of-two grouping 2. Under supported D256, padded KV, compiled Volta MMA and applicable layout conditions:

   Q.ne[1]	Selected family	Q8_0 handling
   1	VEC	Direct Q8_0 if the template was compiled
   2-8	TILE	Full passed-view FP16 materialization
   9+	MMA_F16	Full passed-view FP16 materialization

   If the vector Q8_0/Q8_0 template was not compiled, even VEC can convert to FP16.

   The extra-data and launch helpers allocate ggml_nelements(K)*sizeof(half) and the corresponding V buffer, then actually convert the tensor views. Mask-bound processing comes after conversion. Existing split/stream-K, FP32 combination, and mask scanning must be reused rather than presented as missing features.

   A passed view is not necessarily one slot's live length. Trace the actual ne[], nb[], view identity and mask. Do not multiply a unified cache span by four simply because four slots exist.

   For one distinct K/V stream, Hkv4/D256, separate K and V:

   Q8 bytes       = 2 * N * 4 * 256 * 34/32
   FP16 temporary = 2 * N * 4 * 256 * 2
   

   Passed KV span	Persistent Q8 K+V	Full FP16 temporary	FP16 write + one subsequent read
   32768	68 MiB	128 MiB	256 MiB
   131072	272 MiB	512 MiB	1024 MiB
   262144	544 MiB	1024 MiB	2048 MiB

   The last column is logical traffic, not a measured DRAM counter or extra simultaneous allocation. Summing layer writes does not give peak VRAM.

   A proposed W3/Hq24/D256, 32-partition FP32 output/max/denominator layout is 3*24*(256+2)*32*4 = 2,377,728 bytes, about 2.268 MiB. That counts partials only, not all workspace. Partition count 32 is semantically tested, not selected as fastest.

   3. Rejected shortcut: simply force W3 onto VEC

   The existing vector wrapper already supports multiple queries using two-query blocks and tail checks. It is a valuable low-effort native control, but not automatically the winner.

   For eligible W3, current TILE sizing gives query tile4/GQA tile2. VEC gives query tile2/GQA tile1. For a near-full prefix at N131072, with no redundant-read cache reuse:

   Arm	Nominal KV passes	Accounted operand traffic
   Stock TILE plus conversion	3 FP16 consumer passes + conversion	2320 MiB
   Existing Q8 VEC forced as a W3 control	12 Q8 passes	3264 MiB
   Proposed grouped direct-Q8 W3/GQA6	1 Q8 pass	272 MiB, excluding partial traffic

   These are output-tile pass counts, not instruction or HBM traces. Partitions divide KV work; they do not each scan the full span.

   Forcing VEC can remove the FP16 buffer yet increase total reads. In the toy cache model with the same redundant-read reuse fraction for both arms, its byte comparison flips near 48% reuse; this is not a measured cache-hit rate.

   Even the conservative one-consumer-pass comparison gives (272+512+512)/272 = 4.765 in logical bytes. Neither 4.765 nor the nominal 8.529 byte ratio is an attention or TG speedup. The sensitivity grid includes candidates losing because math, decode, bandwidth efficiency or other costs dominate.

   The first GPU comparison should therefore be stock TILE vs existing Q8 VEC W3 control vs new grouped Q8 W3, not a global dispatcher override.

   4. PP: a tiny TG tile can regress badly

   For eligible W1024, the inspected SM70 MMA route uses query tile32/GQA tile2, giving ceil(1024/32)*ceil(6/2)=96 nominal KV passes. A proposed all-six-head Q8 tile trades query reuse for head reuse: grouping six heads does not independently confer a 3x gain when fewer queries fit the same matrix tile.

   At N131072 with no modeled redundant-read cache reuse, the baseline conversion plus 96 FP16 passes accounts for 49936 MiB:

   Proposed grouped-Q8 query tile	Q8 passes	Accounted Q8 traffic	Baseline/candidate byte ratio
   4	256	69632 MiB	0.717: worse traffic
   8	128	34816 MiB	1.434
   10	103	28016 MiB	1.782
   16	64	17408 MiB	2.869

   These are cumulative accesses, not VRAM capacity. This rejects “the W3 kernel is also the best prefill kernel.” PP tile10/tile16 are better traffic candidates, but useful math rate and resource limits still determine whether either actually wins.

   NVIDIA's Volta guide gives 96 KiB shared memory per SM and requires explicit dynamic opt-in above 48 KiB. For the proposed W3/GQA6 layout, 18 useful rows padded to 24:

   • KV tile32, single buffer: minimum 44 KiB, at most two CTAs from shared memory alone.
   • KV tile32, double buffer: minimum 76 KiB, at most one CTA, requires opt-in.
   • KV tile64, double buffer: minimum 140 KiB, rejected before metadata.

   These are minimum budgets, not ptxas occupancy results. Register allocation, additional arrays and padding can worsen them. Do not increase MTP depth merely to fill a convenient MMA tile.

   5. GEMM: dequantization fusion is not universally faster

   The native quant layouts use 176 bytes/256 weights for Q5_K and 210/256 for Q6_K. Optimistic Q5_K weight-only arithmetic intensity is about 2.91 FLOP/B at M1, 8.73 at M3, versus 2979 at M1024, excluding activation/output traffic, repeated loads and integer work. NVIDIA's GEMM guide explains why these regimes need different tile/reuse tradeoffs and why roofline alone is insufficient.

   Counterexample in the simulator, weight traffic only:

   M1024, Q5_K, hypothetical baseline token tile128, no redundant-read cache reuse
   full-dequant baseline: 0.6875 + 2 + 8*2 = 18.6875 bytes/weight
   fused tile32:          32*0.6875 = 22 bytes/weight
   fused tile16:          64*0.6875 = 44 bytes/weight
   

   The full FP16 temporary disappears, but packed weights and decode work repeat too often. Larger tiles and sufficient useful math throughput can win; merely deleting the temporary does not prove it. Baseline tile choices and the simulator's illustrative 600 GB/s / 60 TFLOP/s parameters are assumptions, not host measurements.

   The current MMVQ already handles multiple ncols_dst. Do not compare new weight reuse against an invented sequence of M independent GEMVs. Repeated source helper calls also do not prove M-fold DRAM traffic after compiler/cache reuse.

   6. Why the PP/TG priority is workload-dependent

   Using the public Qwen configuration as a structural surrogate, the modeled trunk projections total 24,350,556,160 weights. This includes query gating and recurrent projections; it excludes output-head selection, recurrent updates, small ops and MTP. It is not a census of the owner's mixed-quant GGUF.

   The full-attention/trunk-projection FLOP ratio is:

   Case	FLOP ratio, not runtime share
   Cold 8192-token prompt	0.033
   Cold 32768-token prompt	0.132
   Cold 131072-token prompt	0.529
   Cached 131072-prefix + 1024 new tokens	1.062
   Cached 262144-prefix + 1024 new tokens	2.121

   This supports prioritizing attention more strongly for long cached-prefix PP than for short cold PP. It does not establish per-operator time fractions.

   The existing GDN graph already has fused/chunked alternatives. The numerical GDN experiment validates block algebra and selected states, not a new faster native kernel. Profile the actually selected route, preprocessing and intermediate tensors before choosing a GDN change.

   Graph replay is already supported on Volta. Device-resident MTP handoff may remove a critical sync, but payload size alone cannot rank it: W3/hidden5120 has only 120 KiB of hypothetical FP32 round-trip payload per sequence, while the latency of waiting is unknown. PP1024 has 20 MiB one-way payload.

   For fixed MTP depth2, an illustrative conditional-acceptance model gives TG=(1+p+p^2)/(verify_W3+2*draft_step+catchup_and_other_costs). Keep acceptance and draft depth fixed when testing a kernel. A donor's deeper speculation or different NVFP4 artifact is not a causal gain for this GGUF workload.

   An ideal independent-work pipeline simulation with assumed stage costs31/31/30/8 gives about 3.016x at 32 work items, but this is a zero-transfer upper bound, not an available fork speedup. The current engine may already overlap work, and single-sequence dependent TG cannot use that same independent-job assumption. Preserve the selective P2P/NVLink-island contract.

   7. Conditional priority and exact native comparison

   Workload	First native candidate
   TG/MTP2, long KV	Direct-Q8 W3/GQA6 attention; include stock and existing-VEC controls
   PP, long cached prefix	Separately sized Q8-aware D256 attention; qualify tile10/tile16 and useful math rate
   PP, short/medium cold prompt	Actual SM70 GEMM shape/quant dispatch; then measured GDN subgraph gaps
   TG, short context	Q5_K/Q6_K explicit decode/metadata reuse and layout tuning against existing batched MMVQ
   Graphs, hidden-state handoff, multi-GPU	Only after measured critical-path sync/graph/communication gaps

   Required native evidence:

   • Record actual quant histogram, attention shapes/views/masks, W per sequence, current route and template availability, conversions, peak workspace and graph behavior.
   • Test stock/TILE, existing Q8 VEC W3 control and new grouped W3; keep unsupported-path fallback. Use a separate PP tile policy.
   • Validate native activation quantization, precision, masks/tails, sequence/state isolation, accepted-prefix rollback, relevant optional features and memory limits. CPU numerical PASS is insufficient.
   • Inspect compiler register/shared-memory usage and synchronization correctness. Include conversion, partition reduction and setup in operator timing.
   • Alternate control/candidate full PP/TG runs at unchanged GGUF, Q8 KV, MTP2, b2048/ub1024, device order and safe transport. Include cold PP, cached-prefix suffix PP, ordinary W1 TG and real multi-slot MTP shapes.
   • Report complete-model benefit beyond run variability. If an operator is 30% of time and becomes 2x faster, total speedup is only 1/(0.7+0.3/2)=1.176, not 2x. Actual fractions and component gains must be measured.

   No timing model was fitted to issue #5's target-only PP total: that does not identify operator fractions and excludes MTP. No donor headline tok/s was used as a calibration factor. The 4080 sensitivity combinations are not a success probability.

   Reproducible artifact

   Delivered to the owner in the accompanying ChatGPT response: v100-simulation-lab-20261001.zip (23 files, 82359 bytes), SHA-256:

   a93967a04333fab8a28016bd65acdf4fa1e008da48b7312c370c68f1e9615374
   

   It contains the complete independent prototype/simulator source, run.sh, a 3852-word English report, every numerical and sensitivity result, environment/provenance records, reproduction verification and checksums. Run bash run.sh in the extracted v100_sim_lab directory; it writes a separate rerun-results directory and does not contact or alter the server. The archive is not committed to this repository and this comment is not a deployment claim.

   Status: CPU semantic validation PASS; analytical research complete for this iteration; actual V100 performance PARTIAL/unmeasured. Keep issue #9 open.
3. novkien commented on Oct 1, 2026

   

   novkien

   on Oct 1, 2026 – with ChatGPT Codex Connector

   OwnerAuthor

   More actions

   Implementation and native acceptance plan: Qwen 27B, 50k context, direct llama-server

   Owner direction, 2026-10-01: prepare the detailed implementation and speed-validation plan here. Use the existing Qwen 27B model at 50k context. Start llama-server directly and send requests directly to that process; do not use llama-proxy, its launcher, routes, metrics or dashboard.

   This comment is the implementation/test specification, not a report that code has been implemented or that a GPU benchmark has passed. The preceding simulation report establishes candidate mechanisms, not native speedups. The latest 50k test requirement replaces the earlier proposed 128k/256k performance workload for acceptance.

   Rechecked fork master: e58b770f3c2346deab036d9e4c3b2ea61ca0bbc5. PR #4 and PR #8 are merged. Preserve their selective-P2P and MTP behavior. At implementation time, rebase the design assumptions against current source without resetting, switching or rebasing a shared working checkout.

   1. Deliverables and implementation order

   Implement and measure independent changes, then combine only qualified winners:

   Step	Deliverable	Required exit evidence
   A	Direct-server benchmark driver, fixed token fixtures, baseline manifest, opt-in shape/route instrumentation	An unchanged native Qwen run at the requested context; valid PP/TG counters; measured shapes and A/A variability
   B	Stock/W3-vector controls and a direct-Q8_0 SM70 W3 attention candidate	Correct native outputs, bounded workspace, actual route hits and native W3 performance
   C	Separate PP tile/dispatch policy for direct-Q8 attention	Real 1024-query/tail coverage and improved 50k cold or cached-suffix PP, without a TG regression
   D	Shape-specific Q5_K/Q6_K GEMM and/or selected GDN optimization only if the 50k profile identifies remaining dominant work	Individual ablation proving its contribution; no unmeasured blanket repack/fusion
   E	Final combined binary, regression results, benchmark data and maintained documentation	Repeated, unprofiled, direct-server acceptance against the matching baseline

   Steps B and C are the first concrete kernel work. Step D is not permission to rewrite unrelated runtime subsystems. CUDA Graph stabilization, GPU-resident MTP handoff and topology changes remain separate candidates until a trace demonstrates their critical-path cost. Do not force a simulated favorite to win at 50k by changing the model, quantization, batch sizes or speculation depth.

   Proposed exact implementation targets (new paths below are deliverables to create, not claims that they already exist):

   • Modify ggml/src/ggml-cuda/fattn.cu: eligibility, stock controls and measured SM70 dispatch.
   • Modify ggml/src/ggml-cuda/fattn-common.cuh: shared planning/allocation integration where needed; preserve existing mask, split and reduction contracts.
   • Create ggml/src/ggml-cuda/fattn-sm70-q8.cuh and fattn-sm70-q8.cu: local Q8 decoding, Volta operand mapping, online softmax and bounded partition combination. Reuse helpers in mma.cuh and existing quant types instead of duplicating them.
   • Update ggml/src/ggml-cuda/CMakeLists.txt only as necessary for registration/templates; its current CUDA source glob requires reconfiguration after adding files. Preserve HIP/MUSA and other CUDA builds.
   • Extend relevant cases in tests/test-backend-ops.cpp; use existing server test infrastructure for lifecycle/sequence regressions. Do not create a competing general test framework.
   • Create scripts/bench-sm70-qwen27b.py: stdlib-based direct-process/direct-HTTP fixture, repetition, SSE and result runner. No proxy dependency.
   • Create docs/development/sm70-q8-attention.md: dispatch/precision/workspace contract, exact reproduction and results; link it from the appropriate existing documentation.
   • Only after Step D is selected: modify mmvq.cu, the actual matmul selector in ggml-cuda.cu, or the measured GDN path in gated_delta_net.cu / src/models/delta-net-base.cpp. Do not include all of them by default.

   Keep changes reviewable by these boundaries. No production merge, service restart or deployment is part of posting this plan.

   2. Fixed workload and what “50k context” means

   Use the same existing Qwen3.8-27B UD-Q5_K_XL GGUF identified in this issue, not another Qwen 27B quant, a freshly downloaded replacement, or NVFP4. Confirm general.architecture=qwen35, actual attention geometry, tensor-type histogram and MTP availability from the local artifact. Record every GGUF shard and any separate MTP artifact with SHA-256. The filename alone is insufficient.

   For the primary performance campaign:

   Setting	Fixed value
   Context requested from llama-server	50,000 tokens, --ctx-size 50000
   Slots	1 for attribution; a separate four-slot regression is defined below
   Long prompt actually processed	49,152 token IDs, including the chosen template/special tokens
   Output budget	512 tokens; 49,152 + 512 = 49,664, leaving 336 tokens below the requested limit
   Cached-suffix case	Prime 48,128 tokens, then extend to 49,152 with a 1,024-token suffix
   Target K/V	Q8_0 / Q8_0
   GPU placement	Four local V100 16GB, layer split, same verified UUID order and 31,31,30,8 proportions in every arm
   Batch / microbatch	--batch-size 2048 --ubatch-size 1024
   CPU threads	--threads 20 --threads-batch 20
   Speculation	Two separately launched profiles: none, and native draft-mtp, maximum 2 draft tokens
   Sampling for timed runs	Greedy, fixed seed, no grammar or probability export, fixed output budget
   Input type	Text only; no image encoding or projector in these performance cases

   This is a 50k-capacity, near-full-context workload, not a claim of feeding exactly 50,000 prompt tokens and then generating beyond a 50,000-token limit. Record allocator padding and effective context from startup and /props; retain the logical workload limits above. A short prompt with only -c 50000 does not satisfy acceptance.

   No YaRN/RoPE override, context shifting, automatic context reduction, quant change, MTP-depth increase, RPC device, or CPU layer fallback may be introduced to rescue a candidate's result. Check actual target/draft allocations and layer placement, not just argv.

   For the separate four-slot regression, retain --ctx-size 50000 --kv-unified --parallel 4. This is one shared 50k pool, not four 50k reservations. Use four distinct 12,000-token prompts and 128 generated tokens each (48,512 total), plus a ragged case with lengths 8,000/10,000/14,000/16,000. Keep all active sequence budgets within the pool. Do not silently expand to 200k or advertise four simultaneous near-50k sequences.

   3. Establish a trustworthy native baseline

   The runner must collect before benchmarking:

   1. Source revision and applicable existing modifications, compiler/CUDA versions, CMake configuration, binary and loaded-library hashes. Build immutable A and B artifacts without disturbing shared edits. Never rebuild a running arm's libraries in place.
   2. Model/shard/draft hashes and extracted GGUF metadata; actual target and draft device, cache-type and context settings. Embedded MTP is valid only if the model includes the required tensors. A detached draft must use the same verified artifact in both arms.
   3. GPU UUID/PCI mapping, nvidia-smi topo -m, GPU/process/memory snapshots, clocks, power limits and temperatures. Record the exact GGML_CUDA_P2P_PAIRS mapping after CUDA_VISIBLE_DEVICES ordering. Do not assume physical GPU numbers are process ordinals.
   4. Full sanitized argv/environment, native --version, --help, --list-devices, startup log, /health and /props. Verify that the listener belongs to the child process just launched.
   5. A/A runs with identical binaries/settings to measure normal variability before picking a threshold or tuning candidate parameters.

   Use an available GPU window. Do not kill unrelated jobs or invoke proxy APIs to unload/restart services. Insufficient VRAM or interference is a recorded blocked/invalid experiment, not a reason to alter the fixed configuration. This resource check does not require globally serializing unrelated repository work.

   Build baseline and candidate identically. Use a CUDA 12.x toolkit that can build SM70, not the driver-reported “CUDA Version” as a compiler selector. CUDA 13 removed offline Volta compilation/library support N1. Example for each already-resolved source snapshot:

   # SOURCE, BUILD and CUDA12 are verified absolute paths, not historical host paths.
   : "${SOURCE:?}" "${BUILD:?}" "${CUDA12:?}"
   cmake -S "$SOURCE" -B "$BUILD" \
     -DCMAKE_BUILD_TYPE=Release \
     -DCMAKE_CUDA_COMPILER="$CUDA12/bin/nvcc" \
     -DCUDAToolkit_ROOT="$CUDA12" \
     -DCMAKE_CUDA_ARCHITECTURES=70 \
     -DCMAKE_CUDA_FLAGS="-lineinfo -Xptxas=-v" \
     -DGGML_CUDA=ON -DGGML_CUDA_FA=ON -DGGML_CUDA_GRAPHS=ON \
     -DGGML_CUDA_FORCE_MMQ=OFF -DGGML_CUDA_FORCE_CUBLAS=OFF \
     '-DGGML_CUDA_FA_QUANTS=q4_0-q4_0;q8_0-q8_0;f16-f16;bf16-bf16' \
     -DLLAMA_BUILD_TESTS=ON
   cmake --build "$BUILD" --target llama-server test-backend-ops -j "${BUILD_JOBS:-4}"

   These flags follow current build options S4. Freeze the remaining relevant build options equally. Confirm the Q8_0/Q8_0 vector template and native SM70 code are present; otherwise a control can unexpectedly materialize FP16. Record linking so A cannot accidentally load B's libggml-cuda or a production library.

   4. Implement the W3 attention path without changing semantics

   Observed starting point: for the eligible D256/GQA6 shape, the current SM70 selector uses effective GQA grouping 2; W1 takes VEC, W2-W8 take TILE, and larger widths take MMA_F16. TILE/MMA materialize the passed Q8 K/V view as FP16 S1, S2. The old W4 donor specialization cannot simply be renamed W3.

   4.1 Eligibility and a single shared execution plan

   Create one deterministic plan used by support/dispatch, allocation sizing and launch. It must carry the selected route, tile geometry, partition count, required alignment and byte counts. Do not select one route when reserving scratch and a different route at launch.

   The first new fast path is restricted to native NVIDIA SM70, Q F32, K/V Q8_0, Dk=Dv=256, Hq/Hkv=6, W=3, one sequence, supported strides/views and an additive mask. Initially exclude ALiBi/max-bias, softcap, sinks, sparse-gather mode and other unimplemented variants. Decide eligibility from tensor/op properties, never from a model filename. Unsupported inputs use the existing path unchanged.

   Cover W1/W2 and multi-sequence inputs as explicit fallback tests; do not claim that four server slots imply an attention width of 12. Measure actual ne[], nb[], mask dimensions and sequence layout. If the target 50k run does not produce eligible W3 calls, report this and qualify the observed shape instead of changing draft depth.

   Add a new, test/debug selector such as GGML_CUDA_SM70_Q8_ATTN_MODE={stock,vec-w3,tg,pp,auto}. This variable does not exist yet and is part of the implementation:

   • stock: preserve the current route and allocator behavior.
   • vec-w3: use the already-compiled Q8 vector route only for eligible W3 as a control.
   • tg: enable only the new W3 path.
   • pp: enable only the separately qualified PP path.
   • auto: combine only measured eligible winners; all other inputs fall back.

   Parse once before capture. Print the selected mode once. A forced candidate that never executes is not a successful benchmark. Keep route/reason/bytes counters opt-in and aggregated; no per-layer console writes in timed release runs. Do not add a new server public API for this experiment.

   4.2 Kernel and numerical contract

   Load ordinary block_q8_0 data with its FP16 scale and signed int8 values. D256 comprises eight quant blocks. Respect tensor byte strides and the block's actual alignment; a 34-byte block is not automatically 16-byte aligned. Reuse existing safe load/conversion helpers.

   For each KV tile, decode near the consumer and reuse it across the three queries and six related query heads. Use existing Volta MMA building blocks and explicit lane/fragment ownership. Do not use Ampere cp.async, modern-only MMA shapes, BF16 Tensor Cores or FP8 instructions.

   Maintain FP32 running maximum, normalization sum and output accumulation/partition combination. Apply actual additive mask values and causal/sequence restrictions, including padded tails. Handle empty partitions without exp(-inf - -inf) NaNs. Qualify the existing all-masked-row contract before enabling it. Do not drop low-probability terms or change precision/FTZ behavior merely to win a benchmark; any different intermediate rounding must pass native numerical and model checks.

   Start with small single-buffer KV tiles (16/32) and a bounded partition set such as 1/2/4/8/16/32. Measure before adding double buffering. NVIDIA documents 96 KiB shared memory per Volta SM and explicit opt-in for dynamic allocations above 48 KiB N2. Account for all arrays, padding, registers and spills, not only the KV buffer.

   4.3 Workspace, graphs and fallback

   Do not allocate a full FP16 K/V copy for the new route. For a partitioned FP32 result, account for at least W * Hq * (Dv + 2) * partitions * sizeof(float), plus alignment/metadata and actual liveness. Set an initial deterministic per-op scratch cap of 8 MiB for the W3 candidate; it is a proposed engineering budget, not a measured requirement. P=1 may write the normalized output directly where safe.

   Use the existing backend allocator/pool, streams and capture-safe ownership. No lazy allocation, device synchronization, environment lookup or host data-dependent shape decision inside capture. Recompute plan keys on relevant shape/stride changes. Verify graph replay after alternating prompt tails and W1/W3 generation.

   If a shape/resource combination is ineligible, choose stock before allocation/launch. Do not attempt to recover from a poisoned CUDA context after an illegal access or OOM and call it fallback. Keep supports_op, allocator bytes and launch semantics consistent.

   5. Implement PP separately, then profile remaining PP work

   Reuse the correct Q8 load/softmax machinery, not the tiny W3 query tile. Test query tiles 8/10/16 (including their Volta padding and tail masks) against the existing MMA path for the actual 1024-query microbatch and its tails. Six-head packing can reduce the number of queries that fit; it is not a free 3x reuse multiplier.

   Begin PP with no KV partitioning or very few partitions when query tiles already expose enough parallelism. Propose a 64 MiB maximum per-op PP global scratch budget, reducing partitioning or selecting stock when needed. Never materialize a full query_count * KV_length score matrix. Actual peak memory on every card, including draft/recurrent state, is still the acceptance criterion.

   Measure conversion + attention + partition reduction together. A candidate that removes FP16 storage but repeatedly decodes/reads Q8 too often can lose. Final dispatch needs a simple measured shape/length boundary, not a global “always direct Q8” switch or runtime autotuning during a request.

   At 50k, cold PP may still be dominated by projections/GDN. The diagnostic trace must separate full attention, weight conversion/GEMM, recurrent/GDN work, draft catch-up, copies/synchronization and graph overhead. Do not infer shares from FLOPs or layer counts.

   If Step D is warranted:

   • Q5_K/Q6_K: compare against current multi-column MMVQ, including Q8_1 activation conversion. Hoist/reuse quant metadata and weight decoding only where native instructions show duplication. Profile real M/N/K and quant types; do not compare against M independent GEMVs. Preserve FP32 accumulation policy.
   • PP GEMM: include all dequantization/setup/workspace costs. Avoid mandatory weight repacking initially; add a backend representation only after its lifetime, tied weights, fallback, readback and peak-memory contract is defined.
   • GDN: identify fused versus chunked dispatch before changing it. Optimize the measured transforms/GEMMs/solve or recurrence; preserve final state and all accepted-prefix rollback snapshots. Existing register state and chunked graphs are not new features to add.

   Each extra candidate gets its own ablation. If only TG improves, report TG improvement and PP neutral; do not describe the entire PP/TG goal as completed.

   6. Direct llama-server launch contract

   The benchmark driver must start the actual binary with subprocess.Popen(argv, start_new_session=True) and connect only to its loopback listener. Use urllib.request.ProxyHandler({}) or equivalent to prevent environment HTTP proxies. No proxy launch job, routing API, reverse proxy, production port or dashboard scraper participates.

   Below is the native launch form to implement. Resolve paths, UUIDs and the unused port from local preflight; these variables deliberately do not assert stale host facts. PROFILE is none or mtp2, SLOTS=1 for primary results and SLOTS=4 only for the bounded regression above.

   set -euo pipefail
   : "${BUILD:?}" "${CUDA12:?}" "${MODEL:?}" "${V100_UUIDS:?}" "${P2P_PAIRS:?}" "${PORT:?}"
   : "${PROFILE:?}" "${SLOTS:?}" "${ARM_KIND:?}"
   SERVER="$BUILD/bin/llama-server"
   
   COMMON=(
     --model "$MODEL" --alias qwen27b-sm70-50k
     --host 127.0.0.1 --port "$PORT"
     --ctx-size 50000 --parallel "$SLOTS" --kv-unified
     --batch-size 2048 --ubatch-size 1024
     --threads 20 --threads-batch 20
     --device CUDA0,CUDA1,CUDA2,CUDA3
     --split-mode layer --tensor-split 31,31,30,8 --n-gpu-layers all
     --flash-attn on --cache-type-k q8_0 --cache-type-v q8_0
     --fit off --no-context-shift --no-mmproj --no-webui
     --cache-ram 0 --no-cache-idle-slots
     --ctx-checkpoints 32 --checkpoint-min-step 8192
     --warmup --perf --log-colors off --log-timestamps --log-verbosity 3
   )
   
   SPEC=(--spec-type none)
   case "$PROFILE" in
     none) ;;
     mtp2)
       SPEC=(--spec-type draft-mtp --spec-draft-n-max 2)
       # Set these two variables from the verified native MTP fixture, equally for A/B.
       : "${DRAFT_CACHE_K:?}" "${DRAFT_CACHE_V:?}"
       SPEC+=(--spec-draft-type-k "$DRAFT_CACHE_K" --spec-draft-type-v "$DRAFT_CACHE_V")
       # An embedded-MTP artifact needs no separate draft model flag.
       if [[ -n "${DRAFT_MODEL:-}" ]]; then
         SPEC+=(--spec-draft-model "$DRAFT_MODEL")
       fi
       # Preserve verified draft placement when explicitly configured.
       if [[ -n "${DRAFT_DEVICES:-}" ]]; then
         SPEC+=(--spec-draft-device "$DRAFT_DEVICES")
       fi
       ;;
     *) echo "Unsupported PROFILE" >&2; exit 2 ;;
   esac
   
   RUN_ENV=(env -i
     "PATH=$PATH" "HOME=$HOME" "LANG=C.UTF-8"
     "LD_LIBRARY_PATH=$BUILD/bin:$CUDA12/lib64"
     "CUDA_VISIBLE_DEVICES=$V100_UUIDS"
     "GGML_CUDA_P2P_PAIRS=$P2P_PAIRS"
   )
   case "$ARM_KIND" in
     baseline) ;;
     candidate)
       : "${ATTN_MODE:?}"
       RUN_ENV+=("GGML_CUDA_SM70_Q8_ATTN_MODE=$ATTN_MODE")
       ;;
     *) echo "Unsupported ARM_KIND" >&2; exit 2 ;;
   esac
   
   exec "${RUN_ENV[@]}" "$SERVER" "${COMMON[@]}" "${SPEC[@]}"

   The candidate-mode variable above is the new implementation control from section 4, not an existing native CLI option. Stock baseline must not receive an invented feature flag. Verify the candidate's startup marker rather than assuming an ignored environment variable worked.

   Current native syntax is --spec-type draft-mtp --spec-draft-n-max 2; --draft-max is removed. Current device syntax is --device, not an assumed --devices alias S3. Confirm these against the built artifact's help before execution. Resolve default/explicit draft offload, split and cache behavior from actual startup; if extra placement arguments are needed to preserve the verified fixture, freeze and include them in the recorded argv of both arms. Do not silently disable MTP if the artifact cannot initialize it S7.

   The disabled RAM archive/idle-slot policy is an explicit isolated benchmark profile, not a production configuration change. In-memory per-slot prefix reuse and checkpoints remain testable. Keep it identical across A/B and report it.

   Runner lifecycle requirements: reserve/check an unused port, start only one tested arm at a time on the same GPUs, capture stdout/stderr, poll its health until ready while also checking child exit, and use an external monotonic deadline. Suggested defaults are 600 s for readiness and 1800 s per long request, recorded as timeout limits, not performance targets. On error save evidence; do not auto-restart and overwrite the failed trial. Stop only the owned child/process group and wait for it to exit. No pkill llama-server, systemd restart, clock/power-limit change, or proxy operation.

   7. Deterministic fixtures and direct request sequence

   Implement three frozen text fixtures: technical prose, code, and structured records with nontrivial cross-record questions. Do not fill 49k with one repeated token/line that makes MTP artificially easy. Store their generation seed, source text, template choices, token arrays and hashes.

   Use the baseline native /apply-template and /tokenize during preparation, outside timings. Reserve the correct template prefix/suffix and adjust only the body token budget to obtain 49,152 valid IDs. Explicitly include required special tokens once. /completion accepts an integer token array and does not automatically prepend BOS to this form S3. Submit the same arrays to every arm; verify compatible tokenizer metadata.

   Produce these payloads per fixture:

   • cold.json: complete 49,152-token array, cache_prompt=false, 512 output tokens.
   • prime.json: first 48,128 IDs of that same array, cache_prompt=false, n_predict=0.
   • cached.json: complete 49,152-token array, cache_prompt=true, 512 output tokens.

   Use id_slot=0 for primary runs. Common timed-generation fields:

   {
     "n_predict": 512,
     "temperature": 0.0,
     "seed": 424242,
     "samplers": ["temperature"],
     "repeat_penalty": 1.0,
     "presence_penalty": 0.0,
     "frequency_penalty": 0.0,
     "ignore_eos": true,
     "stop": [],
     "n_probs": 0,
     "n_cache_reuse": 0,
     "id_slot": 0,
     "stream": true,
     "return_tokens": true,
     "return_progress": false,
     "timings_per_token": false
   }

   Add the full token-array prompt and the appropriate cache_prompt to make each complete payload. For the prime request set n_predict=0 and stream=false. No generated tokens from a priming completion are needed. Native documentation explicitly supports cache-only evaluation with n_predict=0 S3.

   The native request path, usable directly from a second terminal while that server runs, is:

   : "${RUN:?}" "${PORT:?}"
   mkdir -p "$RUN"
   BASE="http://127.0.0.1:$PORT"
   curl --noproxy '*' --fail --silent --show-error "$BASE/health"
   curl --noproxy '*' --fail --silent --show-error "$BASE/props" > "$RUN/props.json"
   # Cold PP + TG: generated fixture/cold.json already contains the integer prompt.
   curl --noproxy '*' --fail --silent --show-error --no-buffer --max-time 1800 \
     -H 'Content-Type: application/json' --data-binary @fixture/cold.json \
     "$BASE/completion" > "$RUN/cold.sse"
   # Prime is outside the cached-suffix timer, but keep its cost and response.
   curl --noproxy '*' --fail --silent --show-error --max-time 1800 \
     -H 'Content-Type: application/json' --data-binary @fixture/prime.json \
     "$BASE/completion" > "$RUN/prime.json"
   curl --noproxy '*' --fail --silent --show-error --no-buffer --max-time 1800 \
     -H 'Content-Type: application/json' --data-binary @fixture/cached.json \
     "$BASE/completion" > "$RUN/cached.sse"

   The Python runner additionally timestamps request start, every token-bearing SSE event and completion with a monotonic high-resolution clock. Parse SSE records, not TCP reads/lines; ignore heartbeats and empty events. One event can carry several accepted MTP tokens, so event count is not token count. Deduplicate any final response token representation rather than counting it twice. Preserve raw responses for auditing.

   Cache validity is compulsory. For cold runs require timings.cache_n=0 and the expected processed prompt length. For cached runs retain the prime result and inspect timings.cache_n, prompt_n and checkpoint reuse. A recurrent model may re-evaluate a boundary token or rewind to a checkpoint: compare equal reuse/processed-work states between A/B and use actual prompt_n, not an assumed 1024. Report extra re-evaluation separately. If a genuine cached case cannot be produced, mark that cell PARTIAL; do not label a full prefill as suffix PP. Re-prime before every measured cached case; never let the previous 512 generated tokens become an uncontrolled cache state. Do not import saved slots across binaries or assume target/draft persistence is interchangeable.

   8. Native numerical and regression tests

   Extend the existing backend tests with fixed seeds and replayable shape fixtures; retain their operator-specific reference/error criteria. Do not carry over the CPU simulator's tolerances as proof of HMMA equivalence. Record finite checks, maximum absolute error, relative/normalized error and worst-case location; set any additional tolerance before looking at performance results. Never relax a failing tolerance merely to retain a fast tile.

   Minimum coverage:

   Area	Cases
   Primary shape	D256, Hq24/Hkv4, W3, Q8_0/Q8_0, real 50k-derived strides/views
   Width boundaries	W1,2,3,4,7,8,9,10,16,31,32,33,1023,1024 plus actual PP tails
   KV boundaries	Short valid spans, 49,152, near 50k and actual padded allocation span; unaligned/unsupported spans take safe fallback
   Masks and layout	Causal, additive finite biases, partial last tiles, ragged sequences, masked padding, nonzero offsets, supported non-contiguous strides; excluded ALiBi/softcap/sinks/sparse cases use stock
   Partition behavior	Empty partitions, partition tails, P=1 and each compiled selected P; independent output-max-sum reduction checks
   Numerical stress	Zero/near-zero Q, extreme but valid Q8 scales/values, large score differences, concentrated and diffuse attention; no NaN/Inf
   Device behavior	Actual native V100 execution; non-SM70 eligibility rejection and affected build coverage; no unqualified fast path on newer/other backends
   Graph/allocator	Repeated capture/replay, shape changes, exact scratch sizing/alignment, buffer lifetime and plan/launch agreement
   Model/state	MTP acceptance of zero/one/two drafts, rejection and rollback, slot reuse, cancellation followed by a valid request, cold versus cached continuation, bounded four-slot run

   For all-masked rows, first establish the maintained operator's defined behavior; do not impose a new zero-output convention silently. Candidate behavior must match that contract or the case remains ineligible.

   Use a SM70-supporting Compute Sanitizer installation. Run memcheck, then racecheck/synccheck and relevant initcheck on focused native fixtures, with --error-exitcode 1; a zero application exit alone is insufficient N3. Select the operator/backend filters from the built test-backend-ops --help, saving the exact invocation. Separate sanitizer OOM/tool limitations from kernel findings. Sanitizer/profiler runs are never throughput samples.

   On the full 27B model, retain generated token IDs/hashes for greedy A/B. Investigate any divergence at the same prefix with limited teacher-forced/logit checks outside timing, plus task-specific correctness on the three fixture families. Exact greedy parity is strongest evidence; it is not a blanket claim that all valid floating-point kernels must be bitwise identical. If outputs diverge, numerical and task-quality qualification is required before any win is accepted, and changes to speculative acceptance must be disclosed. Do not claim “same semantic work” solely because both responses contain 512 tokens. Ordinary task-quality checks must also run with normal EOS behavior; ignore_eos=true is only the fixed-work timing control.

   9. Instrumentation: show the candidate actually removes the cost

   In short diagnostic runs record target/draft identity, tensor name, dtype, ne[]/nb[], W, GQA grouping, visible/passed KV span, selected route/reason, query/KV tile, partitions, scratch bytes and FP16 materialization bytes. Aggregate repeated shape records. Record enough chronology to separate prefill, draft catch-up and verification.

   For a winning attention route require evidence of its actual launches and absence of the full-view FP16 K/V conversion on those calls. A process-wide allocation remaining for unsupported fallback does not mean the route failed, and disappearance of one allocation alone is not speed proof.

   Use CUDA events/native backend synchronization for complete operator timing, including conversion, attention and reduction. Use Nsight Systems/Compute only in diagnostic runs with an installed version that supports V100. Inspect memory traffic, Tensor-Core activity, occupancy, shared memory, spills and graph/copy gaps where supported. Nsight Compute can replay, flush caches or serialize work; final E2E timing must therefore be unprofiled N4. Do not change device clocks for just one arm.

   10. Performance matrix and correct PP/TG accounting

   Compare these arms during screening, without conflating their effects:

   • A: unmodified baseline from the current pinned source.
   • B0: candidate binary in stock mode; establish that instrumentation/dispatch scaffolding itself is neutral versus A.
   • Bv: vec-w3 control, not automatically a production choice.
   • Bt: W3 direct-Q8 only.
   • Bp: PP direct-Q8 only.
   • Bbest: only independently qualified paths together; add a selected GEMM/GDN delta only with a separate ablation.

   Required primary cases, always with one slot and requested context 50,000:

   Case	Request work	Profiles	Main measurements
   Cold near-50k	49,152 new prompt tokens + 512 generated	none and mtp2	Cold PP, TTFT, TG, complete request time
   Cached near-50k	48,128-token prime, then extend to 49,152 + generate 512	none and mtp2	Actual suffix PP, boundary re-eval, TTFT, TG; prime cost reported separately
   Short-context regression	A fixed shorter fixture under the same 50k-capacity process	none and mtp2	No short-shape regression; not evidence of long-context improvement
   Four-slot regression	Shared-pool workloads from section 2	none and mtp2	Per-slot correctness/rates and aggregate wall-clock throughput

   At the pinned source, server_slot_stats serializes these metrics and excludes the first generated token from TG rate, since it comes from final-prefill logits S5, S6:

   PP_tok_s = 1000 * timings.prompt_n / timings.prompt_ms
   TG_tok_s = 1000 * max(timings.predicted_n - 1, 0) / timings.predicted_ms
   MTP_acceptance = timings.draft_n_accepted / timings.draft_n  # only if denominator > 0
   TTFT_wall = first token-bearing event time - request-start time
   request_wall = final-completion time - request-start time
   

   Use final timings as the primary native stage metrics; recompute the rates to audit the reported fields. cache_n is reused prompt work and must not enter the PP numerator. predicted_n-1 is emitted-token work for this metric, not the number of verifier launches. Obtain verification-round counts from diagnostic stats/traces rather than assuming they equal token count. Missing or zero timing denominators invalidate the corresponding measurement; no substitute 0/1 values.

   For MTP, report accepted/rejected drafts and observed W1/W2/W3 distribution, not merely the configured depth. Never count draft tokens as user-visible TG output. Compare none A/B separately from mtp2 A/B; comparing candidate MTP-on with baseline MTP-off would measure different features.

   Also report wall-clock token delivery. SSE batching makes per-token GPU latency unobservable from event spacing alone: retain token counts per event, and do not divide one network interval into invented individual-token times. For four slots compute aggregate rate on a common wall interval; do not add rates with unrelated time denominators.

   11. Repetition, selection and explicit success criteria

   First screen candidates using focused native operator tests and three alternating A/B pairs on one frozen long fixture. Freeze the candidate geometry/dispatch before the final campaign. Keep screening data distinct from final acceptance; do not choose a tile on the same trials presented as independent confirmation.

   For each claimed primary workload/profile, run at least five paired repetitions per frozen fixture across the three fixtures. Counterbalance A/B launch order, keep equal warmup, and warm the real PP/W3 shapes outside recorded samples. A clean startup per arm/pair is preferred for controlled cache/allocator state; model loading is outside request timings but is still recorded. Never treat a previous-request prefix as “warmup only” and then count its cache hit as cold PP.

   Record environmental conditions throughout, using the same low-frequency telemetry in both arms. Predeclare invalidation causes: wrong hashes/route/context, unequal processed or cached work, unrelated GPU interference, timeout, truncation/context shift, missing final response, OOM/Xid, or sampler/output-budget mismatch. Preserve every invalid/failed trial and its reason. Re-run a paired block for a documented external disturbance, not because a valid run was slow.

   Proposed practical acceptance thresholds, fixed before the final run:

   • On the claimed target cell, paired geometric-mean throughput ratio is at least 1.05 (+5%), with a 95% paired confidence interval lower bound above 1.00. Compute paired log ratios and bootstrap paired blocks within each fixture (e.g. 10,000 deterministic resamples), then weight fixtures equally. Preserve medians, dispersion and all raw runs. If A/A variability makes this undecidable, the verdict is PARTIAL, not PASS.
   • No required unaffected PP/TG cell has a median throughput regression greater than 3%. Investigate smaller noisy differences; use fallback for a repeatably losing regime instead of averaging it away.
   • Native stage improvement must agree with the corresponding wall-clock behavior. Report TTFT and complete-request latency separately; a suffix PP win with worse complete time is not an unqualified E2E win.
   • Correctness/quality and state tests pass; no new OOM, illegal access, leak trend or forbidden peer operation. Peak VRAM on each 16GB card includes target/draft weights, Q8 KV, recurrent snapshots, scratch and graph allocations.
   • Evidence proves the candidate route was exercised at near-50k, not only that the server loaded successfully. The final combined binary must be measured again after the last relevant source change.

   Use separate verdict columns: cold PP, cached PP, TG without MTP, TG with MTP2, correctness, memory. A kernel may pass one performance goal while another remains neutral or PARTIAL. Do not claim both PP and TG improved unless the respective cells pass. A proven regression/correctness failure is FAILED; an unsupported fixture, noisy interval, absent route hit or missing hardware evidence is PARTIAL.

   The five/three-percent figures are proposed engineering acceptance criteria, not simulated speed predictions or an asserted standard. Better measured gains are welcome; screenshots, peak tok/s and byte-saving ratios alone are not acceptance.

   12. Publication and issue closure checklist

   Attach or commit, through the normal authorized repository workflow, the exact implementation and a durable results directory containing:

   manifest.json                  # source/binaries/libraries/model hashes, hardware and argv
   fixtures/                      # text, token arrays, request JSON and hashes
   runs.jsonl                     # all valid, invalid, failed and warmup runs, labeled
   raw/<arm>/<fixture>/<trial>/    # native logs, final JSON, SSE with monotonic timestamps
   telemetry/                     # topology, memory, clocks, temperature and process evidence
   diagnostics/                   # shape/route/bytes census, compiler resources, native profiles
   correctness/                   # backend/sanitizer/state/quality results
   summary.csv                    # absolute rates, times, ratios, CI and per-goal verdicts
   REPORT.md                      # what won, what lost, limits, exact rerun commands
   

   The driver must make a disconnected/malformed SSE stream, wrong token budget or missing timing field visible as an invalid trial, not silently produce a summary row marked successful. Add parser tests for multi-token MTP events, empty events, final-token duplication, cache counts, zero denominators and process failure. A normal rerun should require one manifest and the two immutable binary directories, not hand-edited responses.

   • Implementation and support/allocation/launch contracts match; untouched shapes/platforms retain their current route.
   • Native Qwen 27B artifact, Q8 target KV, requested 50k context and fixed batch/device settings are proven for both arms.
   • All performance requests ran against directly started llama-server, with no proxy control/data/measurement path.
   • Baseline and candidate-stock controls are comparable; no stale/mixed CUDA libraries.
   • W3 and PP route hits, scratch bounds and native numerical/state tests are evidenced.
   • Cold and cached work are distinguished using actual processed/reused counts; TG uses the native first-token convention.
   • Paired unprofiled measurements and per-goal PASS/PARTIAL/FAILED verdicts satisfy the predeclared criteria.
   • Final combination and the bounded four-slot/short-shape regressions pass after relevant edits.
   • Losing candidates, invalid runs and remaining limitations are retained rather than removed from the report.
   • Documentation and reproducible artifacts are published. Do not close Research: port V100/SM70 kernel and runtime optimizations into llama.cpp-fork #9 merely because a PR builds, a simulator passes, or one microbenchmark improves. Closure must state the actually completed scope; unresolved PP/TG goals remain explicit.

   Source anchors checked for this plan
4. novkien commented on Oct 3, 2026

   

   novkien

   on Oct 3, 2026

   OwnerAuthor

   More actions

   Work report - approved Flash Next implementation intake

   Owner and scope: Codex /root coordinates the owner's explicit 2026-10-03 implementation request for this issue. The new current-scope section records the accepted five-GPU Flash Next workload, separate optimization PRs, end-to-end delivery and energy policy; it supersedes the older direct Qwen 27B research restrictions.

   Done: re-read the body and all three research comments, verified current master 7e0d6c8, found no competing open PR, and established isolated local codex/flash9-* worktrees. No source optimization or GPU speed result is claimed yet. The current route uses IQ3_XXS, Q8 KV, two slots, b/ub1024, c262144, target devices CUDA0..3 plus RPC0, and MTP maximum3. Actual artifact shapes and running-build identity will be captured before candidate qualification.

   Blocker: proxy :8082 has no resident/active model work but still has the host-power failed fence. Owner notified to use the owner-only Reset alarms control; this task will not invoke it or bypass admission. Independent local source/benchmark preparation proceeds.

   Next: /root owns hardware/lease/proxy operations and benchmark acceptance. Separate local workers inspect attention candidates and prepare the reproducible benchmark harness; no worker independently activates GPU workloads or mutates route/power state. The next report will contain baseline/artifact evidence or a concrete blocker update.
5. novkien commented on Oct 3, 2026

   

   novkien

   on Oct 3, 2026

   OwnerAuthor

   More actions

   Work report - owner-authorized exceptional proxy recovery

   The owner explicitly instructed this Codex task to hard-restart the proxy if the existing alarm blocks the Flash benchmark, while another named Codex task implements llama-proxy ggml-org#384. This later direct instruction authorizes this bounded recovery action; it is not ordinary source deployment or authority to reset alarms through the owner-only API.

   Evidence: acquiring the task's agent idle hold returned HTTP409, host power transition already committed; cannot prevent it; no lease or GPU command started. Proxy still reports host_power failed with the power-confirmation error and no loaded model, active request, resource owner or pending MCP work. Current serving release/startup pointer is adfe3aa7ad0f4924af111e589c998f2e2a69e0ac. Management status has no nonterminal restart job; the last ggml-org#542 job b5ee7a62c59f4bfa93d1d92f8507d32a completed with all participants retired. Compatible managed handoff retains the failed power state and thus does not clear this recovery blocker. The separate ggml-org#384 task is authoring its fix locally, not activating a candidate at inspection.

   Next: preserve pre-recovery evidence, recheck restart ownership and activity under the deployment lock, then restart only the proxy user unit at its current startup pointer. Verify old-process retirement, same source release, :8082 readiness, power/admission state and hold acquire/release. Report this as hard-restart recovery, not a safe restart or ggml-org#384 bug fix. No model/GPU/ComfyUI/host reset is included.
6. novkien commented on Oct 3, 2026

   

   novkien

   on Oct 3, 2026

   OwnerAuthor

   More actions

   Work report - exceptional proxy recovery verified

   The owner's explicitly requested hard restart completed at the existing source/startup pointer adfe3aa7ad0f4924af111e589c998f2e2a69e0ac. Preflight under the deployment lock found no nonterminal safe-restart job, resident model, active request, resource owner or pending MCP/Comfy work. Old supervisor/authority/worker PIDs 1186435/1186423/1186434 retired; new supervisor 1408702 serves :8082. Startup pointer and release are unchanged. Read-back alarms=[], power phase=idle.

   This is bounded hard-restart recovery, not a safe-restart job and not acceptance of the ggml-org#384 source fix. No alarm-reset API, GPU-host power cycle, model/Comfy stop or source deployment was performed. The ggml-org#384 implementation remains with its existing task. The Flash optimization task can now attempt its owned short-lived agent hold and resume artifact/baseline collection. Private pre/post operational evidence is retained locally; no private request contents or credentials are published.
7. novkien commented on Oct 3, 2026

   

   novkien

   on Oct 3, 2026

   OwnerAuthor

   More actions

   Work report - exact Flash artifact and installed-runtime screening

   The current IQ3 target is verified as qwen4exp: 48 trunk blocks, embedding2560, attention D256/Hq24/Hkv2 (GQA12), 512 experts/10 active, full-attention interval4 and native context262144. Both shards and the existing Q4_K_M draft were hashed. The second 26.822GiB shard contains only per_layer_token_embd.weight, an existing CPU lookup table; it is not evidence of trunk-layer fallback. Target quantization is mixed, so optimizations must be selected by actual tensor types rather than the IQ3 filename.

   Artifact SHA256: target shard1 219ea929900dfa9ef091f3aa473fdba6874b65fcb36526d7d851ac9e95856d15; shard2 316b46f3a2dbd68c900f43136ab9449f9dcc3725dfd8c794847c204bc161e113; draft b646ef60eaae2a9ed849e75f15f399629ca22633555e99e809959e95f22a1575.

   The installed native/RPC pair is source98b3e092d34b99222d9144f90e388767f02ec117 with an empty patch, even though physical source HEAD is7e0d6c858. Normal proxy Load succeeded in41.24seconds at unchanged c262144/ub1024/two slots/ts11,10,11,10,7/MTP3 on RPC0. Native logs prove target49/49 and draft50/50 offload. Target graph reservation first failed, then retried without pipeline parallelism; final compute allocations are about3520MiB per target GPU plus1210MiB for the RPC draft. Initial free memory in process-device order is2289/3161/723/1071/601MiB. These are startup facts, not proof of the causal share of PP time.

   Three screening requests through proxy8082 used one frozen structured fixture with exactly8192 newly processed tokens (cache_n0) and512 generated tokens, greedy sampling/MTP unchanged. Native PP:453.98,490.28,491.97tokens/s; native TG:61.00,69.43,71.78tokens/s; wall:27.69,33.14,24.77seconds. These are installed-runtime screening results, not the five-pair/three-fixture acceptance campaign or a claimed optimization gain. The first run's slower warm state and wall/native gap remain explicit.

   Next: matched immutable host/RPC baseline builds from7e0d6c858 are now compiling (CUDA12.9/SM70 native; CUDA13.3/SM120a guest). This separates upstream mask changes from our candidates. Opt-in local diagnostic patches collect actual attention dispatch/shape and MTP getter/catch-up cost before kernel/dataflow changes. An independent local placement analysis uses per-layer GGUF sizes and startup allocations. No optimized PR, speed gain or final context selection is claimed yet. Owned wake leases were released after completed inspection/screening commands; a lease is now renewed only while the actual build commands run.
8. novkien commented on Oct 3, 2026

   

   novkien

   on Oct 3, 2026

   OwnerAuthor

   More actions

   Work report - owner requires pipeline parallelism

   After seeing the installed-runtime allocation fallback, the owner explicitly clarified that disabling pipeline parallelism is not acceptable. Final acceptance now requires the target's pipeline parallelism to remain enabled with no allocation fallback. Existing pipeline-off screening is retained only as historical installed-runtime evidence, not a valid final configuration or the pipeline-enabled A/B baseline.

   The context search is constrained by pipeline-on operation. First rebalance actual layer allocations and draft placement, compare ub512/1024 at equal pipeline-enabled settings, and establish the largest stable shared context that meets the owner's PP tradeoff. Retain proof of the final runtime scheduler state, not just the early 'enabled' startup message. No automatic runtime guard or removal of native fallback is requested.

   The unmodified 7e0d6c8 native/RPC baseline build has completed. Next: package matched immutable artifacts, activate through normal model lifecycle, verify pipeline behavior and run comparable measurements. Source diagnostics and offline placement work remain local; no GPU benchmark is run concurrently with a build.
9. novkien commented on Oct 3, 2026

   

   novkien

   on Oct 3, 2026

   OwnerAuthor

   More actions

   Work report - controlled upstream comparison and startup prerequisite

   The owner clarified that real upstream speed gains may justify higher compute VRAM, while pipeline-on remains mandatory. We will preserve and measure those changes, not revert them based on allocation size alone. Upstream ggml-org#29751/66e0c17ee rewrites QSA/indexer pooling; ggml-org#29824/4e2713c16 changes mask construction and notes the compute-buffer cost. Their published speed claims concern other hardware and are not results for this V100/RPC host.

   The new immutable 7e0 native/RPC pair completed its build and was installed through normal proxy unload and complete-package atomic exchange, retaining the old pair. A placement chosen from exact layer sizes (ts12,8,10,9,10; draftCUDA1; c262144; ub512; b1024; two slots; ngl99) successfully reserves target pipeline copies4 without fallback. It then aborts during the draft's startup decode check. The stack and exact reproduction are recorded in prerequisite issue #11.

   A controlled reinstallation of the retained 98b3 pair with that same configuration loads successfully, keeps target pipeline copies4, and serves three exact8192+512 cold SSE requests. All three benchmark records are valid (cache_n0): PP404.71/415.77/415.58, TG66.68/73.69/73.78tokens/s. This separates the 7e startup regression from placement/OOM. It is configuration screening, not final multi-fixture A/B acceptance or an isolated ub tradeoff because the earlier installed screening used a different placement and pipeline state.

   Matching startup graph allocations differ materially: 98b3 uses target compute2969/3085/3085/3085/3089MiB; 7e0 uses4315/4426/4426/4426/4435MiB. Both retain all target layers on GPU. The exact causal/lifetime attribution of this ~1.34GiB/device growth is under local source review; it is not assumed to be all attention conversion.

   Confirmed #11 cause: global trunk QSA metadata causes the MTP graph to register pool inputs although its actual layer48 compress ratio is0. That dense layer never consumes k_idxs, leaving its buffer unallocated while set_input writes it. The minimal compatibility repair constructs the pool only for a QSA-enabled MTP layer and preserves the upstream target graph. Codex /root/flash_attention owns that isolated prerequisite source change under /root coordination; independent review, matched rebuild and live startup/MTP acceptance follow. No optimized kernel or speedup is claimed yet.

   Separately, a newly observed explicit systemd logged-in-user refusal is recorded in llama-proxy#547. The owner-authorized second hard restart restored proxy admission at unchanged ea35908 with old processes retired; no alarm reset API or proxy source patch was used here. Proxy ggml-org#547 is outside this fork implementation and remains open.
10. novkien commented on Oct 3, 2026

    

    novkien

    on Oct 3, 2026

    OwnerAuthor

    More actions

    Work report - benchmark tooling PR opened

    PR #13: #13
    Reviewed head: ae14eb10fc26138f54e7c66b473ad15ace533598.

    This support PR adds the Flash Next proxy stream runner, focused parser/provenance tests, and its usage guide. It records token/cache counts, native PP/TG/TTFT timings, raw SSE artifacts, and immutable host/RPC/model evidence. Missing evidence yields a screening-only result; an environment change invalidates the campaign. This PR adds no inference optimization and makes no speedup claim.

    Validation on the reviewed files: 13 focused unit tests passed, Python compilation passed, and git diff --check passed. Live screening results remain configuration evidence only and do not meet issue #9's multi-fixture paired acceptance.

    At handoff, the issue backlink is present. The PR's Lint, editorconfig, labeler, model-naming, and two Python type-check jobs are still pending; local scoped checks passed and no required branch gate is configured. The optimization issue remains open for the separate source PRs, full A/B results, final pipeline-on context selection, and proxy acceptance.
11. novkien commented on Oct 3, 2026

    

    novkien

    on Oct 3, 2026

    OwnerAuthor

    More actions

    Work report - prerequisite publication, validated runner and first dispatch trace

    Prerequisite PR #12 for #11 is merged at 8d75def06db97c8736989c795be403faf214a9d4. Its maintained tiny Qwen4Exp fixture aborts on unmodified7e0 at the same unused pool-input buffer assertion and passes all3 focused tests after the one-line repair (ordinary architecture, dense MTP, QSA MTP). Exact merged-pair deployment acceptance remains open; the currently running screening pair is still7e0 plus the reviewed correctness patch.

    Benchmark support PR #13 preserves raw SSE bytes, exact token/cache counts, campaign/arm/lane/prompt identities, native/proxy identities and immutable host/RPC binary/library/model hash evidence. Thirteen focused tests and scoped ty0.0.84 checks pass. Incomplete provenance is explicitly screening-only; a mid-run environment change invalidates the campaign. This is supporting measurement work, not a performance optimization claim.

    The temporary diagnostic pair completed an8192+512 request and was removed through normal unload/package exchange afterward. Actual native traces show SM70, Q8_0 K/V, D256,24 query/2 KV heads (GQA12): decode/MTP widths1/2/4 take the F16-converting tile path, PP512 takes MMA-F16. At roughly8.5k-9k KV cells, each K/V F16 scratch is about8.7-9.2MB; those are allocation sizes, not measured traffic. RPC uses direct-Q8 vec for width2 and MMA-F16 for width4/PP. The first per-op profiler overflowed its512-key limit and omitted fused groups/RPC teardown totals, so its operation ranking is incomplete. The revised temporary diagnostic expands bounded keys, measures fusion groups and emits cumulative snapshots; profiled timings will not enter speedup acceptance.

    Source and process evidence now confirm native graphs disabled by environment presence and fusion enabled (GGML_CUDA_DISABLE_FUSION=0); RPC does not inherit native environment and currently uses graph captures. A route value0 cannot enable native graphs because that switch uses presence semantics and the launcher always inserts it. No proxy runtime change is authorized or made. Per-row MTP getters measured62.938ms across8803rows under the intrusive profile, too small there to justify a speedup PR; bulk waiting remains required. QSA score in-place addition is already eligible for gallocr reuse, so no redundant in-place-only PR is planned.

    Configuration screening has now proved ub1024 can fit with pipeline ON: c65536, b1024, two slots, ts12,8,10,9,10, draftCUDA1,49/49 target and50/50 draft GPU layers. Target compute buffers are2276/2513/2513/2513/2516MiB, well below the c262k estimate. Three exact cold8192+512 requests pass: PP484.63/487.93/487.17 and TG78.07/79.53/79.83tokens/s. These are single-fixture screening results, not five-pair acceptance. A fixed-scratch extrapolation from c262k to smaller contexts would be invalid because QSA scratch scales with context and ubatch; the same-context ub512 arm is next. The final largest stable context and near-full-pool generation remain unverified.

    Next source experiment: /root/flash_attention owns a separate family1 direct-Q8 TG probe, starting with existing vector kernels on SM70/D256/GQA12 widths1-4, actual strides/masks, CPU-reference correctness and synchronized microbenchmarks before any live selector claim. Other families await complete cost attribution. Root alone coordinates GPU operations and releases each owned idle lease at the end of the actual command/build/benchmark. No final configuration or1kt/s result is claimed.
12. novkien commented on Oct 3, 2026

    

    novkien

    on Oct 3, 2026

    OwnerAuthor

    More actions

    Work report - PR #13 Lint result

    PR #13 remains open at reviewed head ae14eb10fc26138f54e7c66b473ad15ace533598. Its Lint workflow failed with two findings: NP100 forbids print() at scripts/bench_flash_next.py:1475, and E402 flags the post-sys.path.insert import at tests/test_bench_flash_next.py:10. Labeler and model-naming passed; editorconfig and the two Python type-check jobs were pending. The focused 13-test suite, Python compilation, and git diff --check passed locally.

    No merge was performed. The source owner has assigned these two minimal lint corrections; root review and a new head authorization are required before publication continues. Issue #9 remains open for the optimization and proxy acceptance work.
13. novkien commented on Oct 3, 2026

    

    novkien

    on Oct 3, 2026

    OwnerAuthor

    More actions

    Work report - benchmark support validation update

    PR #13 now points to reviewed head 1ebe411a30851671e59d616fa20d16a66a3bb670. The two CI Lint findings from the prior head are corrected; CI Lint passed on this head (run 37094369075). Editorconfig, labeler and model-naming passed. Two Python type-check jobs remain pending; scoped local type-check, 13 focused tests, Python compilation and git diff --check passed.

    The change remains measurement support and makes no speedup claim. #9 remains open for optimization candidates, paired performance acceptance, final pipeline-on context selection and proxy verification. The support PR is ready to merge at this exact reviewed head; outstanding checks are optional and no required branch gate is configured.
14. novkien commented on Oct 3, 2026

    

    novkien

    on Oct 3, 2026

    Last edited by novkien

    OwnerAuthor

    More actions

    Owner decision - final context must be at least 200k

    The owner clarified directly in the active Codex task on2026-10-03 that the final context must be at least200k tokens. In the already approved unified-KV contract, this is a minimum200,000-token shared -c pool across the two slots. Smaller contexts remain useful controlled comparison fixtures but cannot be the final configuration. The existing pipeline-on/full-GPU-offload requirements and ub512 median-PP-loss<=5% rule still apply; no relaxation was authorized.

    The completed c65536 screen compared the same source, placement, draft, sampling and exact8192+512 fixture: medianPP487.167 with ub1024 versus406.114 with ub512 (about16.64% loss). It is an initial three-run/single-fixture screen, not the final paired campaign. Observed ub1024 VRAM peaks there were roughly12.3-12.9GiB/device (12,608-13,176MiB), and both KV and QSA scratch grow with context. Next: measure ub1024 at c200704 with pipeline/full offload checks, then use measured memory and local source candidates to satisfy the new minimum and maximize capacity. The final near-full-pool test remains required.

    The body now includes this minimum and updates the obsolete initial power-blocker paragraph to the already recorded recovery/current state. No new proxy-runtime authority is inferred.
15. novkien

    added a commit that references this issue on Oct 3, 2026

    bench : add Flash Next proxy stream benchmark harness (#9) [no release]

    ...

    e3b5468
16. 64 remaining items

    Load more

17. novkien commented on Oct 4, 2026

    

    novkien

    on Oct 4, 2026

    OwnerAuthor

    More actions

    Shared resource released; Flash routeproof resumed

    The separate archive stage completed and its process exited. Its own helper receipt records the lease released with final inactive status; a fresh proxy read agreed. Root waited for this evidence before any further power action. Graceful shutdown then completed at 09:39:50 UTC with offline plus fresh 4 W, no hold and no active LLM work.

    The next Flash-only verification job has now passed preflight and started baseline identity setup. It reuses only the complete, hash-bound 100-row CPU and Qwen diagnostic stages; the prior partial Flash group remains excluded. The resumer and bounded load-status read recovery passed root and independent review. The Pi audit also reconciled the earlier accepted Flash load job to succeeded, with slot admission/completion at 09:07:02 UTC; the missing evidence is the diagnostic request, not proof of a model-load failure.

    Both-model TG qualification remains the next phase after Flash routeproof and restored/off cleanup. The prepared final quality/context coordinator has also passed offline root and independent review but remains disabled until the model gain gates justify that acceptance stage. No new production source, route configuration, proxy repair, foreign hold mutation or additional experimental family was introduced.
18. novkien commented on Oct 4, 2026

    

    novkien

    on Oct 4, 2026

    OwnerAuthor

    More actions

    Combined Q6 correctness and actual-route verification complete

    Root completed the missing Flash Next diagnostic request and restored both installed packages to the unchanged 5423 baseline. Graceful shutdown was verified at 09:46:51 UTC, offline with a fresh 6 W reading and no idle hold.

    The combined pure candidate 5423:51b9b224... retains the previously verified 100 passing CPU-reference rows (25 cases x baseline/candidate x native/RPC). The separately built diagnostic candidate 5423:3e0a14e5... now has complete source-bound cold8192+512 structured probes on BOTH Qwen27B and Flash Next. These diagnostic timings are excluded from performance qualification.

    For Flash, the exact benchmark manifest window was 09:44:17.563308-09:44:45.405043 UTC. Its native log contains 31 W1 head branch markers (K2560/M248320, CUDA1/SM70) and two periodic W4 markers (calls4096/8192, K6144/M2560, CUDA0/SM70) inside that request. Startup markers were excluded. The response passed transport/work counters: 8192 processed, zero cache reuse, 512 output tokens, HTTP200 and no sample-invalid reasons. The counter marks branch selection before launch, not individual kernel completion. Qwen's previously completed actual-request W4 proof and CPU references were reused only through pinned completed artifacts; the old partial Flash attempt is excluded.

    Root and an independent reviewer checked the actual request/window and package evidence. No model configuration or production source has changed. Next is the reviewed pure-candidate campaign: five alternating baseline/candidate pairs for each model and each of prose, code and structured cold8k+512 fixtures (20 load groups, 60 measured requests, 20 excluded warmups). A speed claim requires at least 5% native TG gain, a 95% paired interval above parity, and no required case below 0.97. Full quality/context acceptance and PR delivery follow only if those model gates pass; otherwise the baseline remains selected.

    Local evidence: jobs/quant-verify-resume02/result.json SHA256 b35f458b96f5191bed833b17c138e5198c0ed2495d705b1301f74f7818a22285, profile-job/result.json SHA256 3964d89435e7bb52ab5f7751bed497bfda4267d911e8813a76ab8e5b89b7ae9e, and root-routeproof-review.json under the retained flash9-next campaign artifacts.
19. novkien commented on Oct 4, 2026

    

    novkien

    on Oct 4, 2026

    OwnerAuthor

    More actions

    Pure Q6 TG campaign interrupted by a telemetry failure

    The first campaign stopped after seven complete groups. The eighth group (pair1 Flash baseline) completed all three frozen requests, but one of 184 RPC VRAM collector rows records an 8-second SSH timeout inside the code request window. The strict collector-validity gate stopped the campaign; it is not a model/kernel failure or a valid negative speed result. Root has not computed candidate/baseline speed comparisons from this partial run. All original data remains unchanged.

    Both installed packages were restored to exact 5423. The first shutdown attempt was blocked by a foreign agent idle hold. Read-only evidence identified a concurrent archive stage, whose lease interval was 10:17:51.171288-10:25:12.926435 UTC. The stage completed, exited and released its own lease; root never modified that lease or interrupted the task. The existing graceful shutdown helper then completed at 10:26:44 UTC with offline plus fresh 6 W and no hold. Temporal overlap does not prove the archive caused the SSH failure.

    The first six groups (both models' pair0 and both Qwen arms of pair1) completed before the foreign stage and have no collector error rows. Pair1 Flash Q6 code/structured requests overlap that stage; the subsequent Flash baseline warmup and all three measured requests also overlap, including the RPC timeout. To preserve complete paired comparisons and avoid result-based sample selection, the proposed recovery retains only the first six complete groups, excludes BOTH original Flash pair1 arms from qualification with the reason recorded, and reruns that complete Flash pair plus pairs2-4. The excluded raw groups remain available. Independent review is checking the exact retained boundary and evidence validity before the resumer is enabled.

    The resumer will use the same pure source, fixtures, sampling, configurations, five-pair alternating order, telemetry assertions and unchanged analyzer. It will pin the original receipt and retained artifacts, perform fresh baseline/package/model identity setup, run fourteen new groups, and restore/shutdown again. No production candidate has been selected, no performance claim made, and no proxy runtime/configuration or idle-inhibition change is included.

    Original failed/restored receipt SHA256: 30428d2e78038a5680b409fdd95a836150b0063b5698c94e0c5651b264625b6b. Local evidence is retained under flash9-next-20261004/jobs/quant-tg05 and its quant-tg05-power sibling, including the sanitized foreign-stage completion and final shutdown receipts.
20. novkien commented on Oct 4, 2026

    

    novkien

    on Oct 4, 2026

    OwnerAuthor

    More actions

    Reviewed fourteen-group TG resumption started

    Root and independent review validated the retained first six groups before any speed comparisons: all 18 measured runs pass the unchanged request adapter, every raw SSE/prompt/source binding matches, and all three VRAM/slot streams are error-free and cover their request windows. The last retained request ended at 10:14:06.561 UTC, before the concurrent archive interval. Both original Flash pair1 arms remain excluded as a complete pair, with their raw data preserved.

    The separate resumer pins the failed/restored original receipt, the exact six-group records, and all 394 files under those group directories. It passed root full review, independent code review, default/explicit five-pair CLI tests, and offline scheduling/cleanup stubs: six reused groups followed by fourteen fresh groups in canonical order; fresh baseline setup before the first new Flash-Q6 load; cleanup on injected install/request failures. Two local-only wrapper defects (module-global evidence binding and explicit forwarding of the five-pair default) were corrected before any live resumption. The frozen parent, request adapter, pure candidate, telemetry assertions and final analyzer are unchanged.

    The resumer is now active from an offline/no-hold preflight and has started fresh baseline identity setup. Enabled resumer SHA256: 6e1d5faa2b85d23f7069930ffb8b3cdd931251ce638896011bdd68af38169c56; retained-file manifest SHA256: 1a467e6a97ba0988c5cf76de5b4fae2e344be6eeb5ab7ea507d44743c6a0750c. New output is jobs/quant-tg05-resume01/result.json; its eventual composite evidence references the unchanged original six group directories and fourteen new ones. It must finish complete/restored, then verify shutdown, before the unchanged five-pair analyzer runs. No model-level performance claim or new production selection has been made.
21. novkien commented on Oct 4, 2026

    

    novkien

    on Oct 4, 2026

    OwnerAuthor

    More actions

    Second TG interruption and bounded paired recovery

    The reviewed resumer restored exact 5423 after the first new pair2 Qwen baseline group failed the unchanged RPC telemetry gate. All three model requests completed, but the RPC collector recorded seven SSH errors (connection, banner and command timeouts); native GPU and slot streams had no errors. This is invalid telemetry, not a negative Q6 speed result. No comparative speed analysis has been performed.

    A separate archive stage held the GPU host from 10:56:45.533247 to 11:04:16.587021 UTC. It overlaps 18.79 seconds of the newly measured Flash pair1 baseline code window and the entire 27.30-second structured window; all pair2 Qwen baseline windows overlap it. This temporal overlap does not prove the cause of the RPC errors. Both newly measured Flash pair1 arms are excluded as a complete paired unit, and the incomplete pair2 is excluded. The original six validated groups remain the only retained prefix. Every original and resumed raw artifact is preserved.

    The benchmark unloaded its model and restored both installed packages. The first shutdown was blocked by the foreign hold; root did not modify it or interrupt the stage. Once the stage exited and released its hold, the unchanged graceful shutdown helper verified offline with fresh 5 W at 11:05:23 UTC. Failed/restored resumer receipt SHA-256: 363bf58437dee713744d10f6cb1e0530b52123908f031d08b2baaeef59f59013. Recovery boundary: 1185d93c137dcf5100801ba38ad9dfee340094d33ee5e5005f843ee7d3551dd2.

    Root is preparing a bounded recovery that executes one complete model/pair unit per owned job, with baseline setup, two unchanged alternating arms, restoration and verified shutdown. Subsequent jobs must hash-validate the accepted prefix and prior restored segment before reuse; only the final complete 20-group result is eligible for the unchanged five-pair analyzer. A read-only public hold observer will reject an entire segment on foreign hold or observation failure. No hold mutation, relaxed collector assertion, sampling/configuration change, proxy runtime repair or speed-based sample selection is included. The coordinator remains disabled until root and independent review finish.

    Remaining: collect the fourteen valid groups, analyze all six model/family cells, and perform applicable acceptance plus exact-merge delivery only if Q6 qualifies. The accepted production pair remains 5423; no new candidate is selected.
22. novkien commented on Oct 4, 2026

    

    novkien

    on Oct 4, 2026

    OwnerAuthor

    More actions

    Owner stopped competing work; resume the unchanged reviewed campaign

    The owner stated in the current Codex task that the Ren'Py task has been stopped and jarvis-llm is now exclusively available to this work. Its last archive stage exited and released its hold. The host was online again during that stage after an earlier verified shutdown; root corrected its stale status wording. The latest graceful shutdown completed at 11:37:17 UTC with offline, fresh 5 W and no hold, followed by a separate non-multiplexed SSH check.

    This resolves the recurring competing-work dependency. The newly prepared segmented controller remains disabled and will not be used. Root will resume with the already reviewed, unchanged fourteen-group controller (6e1d5faa2b85d23f7069930ffb8b3cdd931251ce638896011bdd68af38169c56) into a fresh jobs/quant-tg05-resume02 directory. Its offline preflight revalidated the original six-group prefix and pinned files. Both interrupted attempts remain excluded exactly as recorded; no comparative speed analysis has occurred.

    The same pure Q6 package, baseline, alternating order, three cold8192+512 fixtures, telemetry assertions and five-pair analyzer remain in force. The sole owned job must restore exact 5423 and verify shutdown before analysis. No idle inhibition is allowed. If it qualifies, the existing acceptance coordinator has a reviewed disabled v2 adding the originally required one-slot near-full check alongside the existing two-slot stage; no new fixed long-cache threshold was added.

    Remaining: fourteen valid groups, all-six-cell TG analysis, applicable acceptance, and exact-merge source/package delivery only for a qualified candidate. Issue #9 remains open and root remains accountable.
23. novkien commented on Oct 4, 2026

    

    novkien

    on Oct 4, 2026

    OwnerAuthor

    More actions

    Work report - owner-requested simpler comparison; Q6 not selected

    Codex /root remains accountable. During the resumed campaign, the owner asked why so many comparisons were needed and requested a simpler approach. Root stopped after the current complete pair, before examining speed results, rather than running the remaining three paired blocks. This changes the measurement plan; it does not create a new confidence-backed gain claim.

    Eight complete groups provide two alternating baseline/Q6 pairs for each model, with 24 measured cold8192+512 requests and eight separate warmups. The six original valid groups were reused with frozen file hashes; the two fresh Flash groups passed the same source, request and raw telemetry checks. Earlier contaminated/invalid pairs remain excluded. The ninth group was interrupted during its pre-load configuration read and produced no measurements. The receipt's failed status records this intentional KeyboardInterrupt, not a model or telemetry failure.

    Model / fixture	Baseline median TG (token/s)	Q6 median TG (token/s)	Paired geometric change
    Qwen 27B / prose	52.44	54.44	+3.81%
    Qwen 27B / code	41.42	42.97	+3.75%
    Qwen 27B / structured	54.30	56.21	+3.51%
    Flash Next / prose	63.63	66.56	+4.55%
    Flash Next / code	65.00	64.38	-0.97%
    Flash Next / structured	69.21	69.87	+0.89%

    These are descriptive two-pair results, without confidence intervals or formal qualification. Qwen paired outputs matched; Flash paired outputs differed, and its structured TG ratios ranged from 0.8829 to 1.1528. PP changed by less than 0.5% in every cell. The modest Qwen results and variable Flash results do not justify more repetitions or promotion from this screen. This is a selection decision, not proof of zero Q6 benefit or a failed five-pair gate.

    The exact-5423 native/RPC pair was restored successfully with no forced cleanup. Graceful shutdown was verified offline at 11:50:57 UTC on 2026-10-04 with fresh 5 W (11:50:50 sample), zero active requests and no idle hold. No idle inhibition, proxy restart, alarm reset, route change, new PR or production Q6 deployment occurred. The immutable candidate remains exact 5423 plus pure patch 51b9b224b2ff6d3168b5f6cf91720d5a6018b7c77844bd1cabdb92a018b4f7d5; its passed 100 numerical reference rows and both diagnostic route proofs are separate from these TG results.

    The local descriptive report is jobs/quant-tg05-resume02-two-pair-summary.json, SHA-256 8e18430d9523d4b94ab2f2a1f7ce390d6916c6f733d18d38103bf491ec3ac000. All original receipts and excluded raw files are retained. Remaining: independent raw-result review and final issue-wide disposition. No additional GPU campaign is scheduled; the previously accepted exact-5423 release remains installed.
24. novkien commented on Oct 4, 2026

    

    novkien

    on Oct 4, 2026

    OwnerAuthor

    More actions

    Final handoff - bounded extension complete, exact5423 retained

    Codex /root reviewed the independent result: all eight completed groups, 24 measured samples and eight excluded warmups are valid; original six-group provenance, raw SSE hashes, native/RPC source identities, unchanged model/startup settings, normal unload and complete error-free telemetry agree with the root analysis. The independent six-cell TG calculation matches the descriptive result. Its local report SHA-256 is 350965208288684fd138c5d3c8c8151ace156db18cd586d3dde696316a7259fc.

    The owner asked to simplify repeated comparisons. The completed two-pair screen gives modest Qwen gains below the original 5% target and variable Flash results; it does not justify another GPU campaign or Q6 promotion. No confidence-backed speedup, five-pair failure, semantic equivalence or zero-effect claim is made. The unresolved statistical variability is a limit of the reported screen, not pending acceptance for a release being installed.

    The Qwen pre-/post-code and MTP2/MTP3 comparison is complete with its five-block report; all considered follow-up candidates have scoped dispositions in the current body. No new source candidate is retained for delivery. Under the original instruction to record unsuitable/no-benefit directions without empty PRs, this concludes the bounded extension. The previously accepted PRs #17, #18, #21 and #22 and exact-merge native/RPC release 5423bddf352adc41a8cc59ed02d38a511b333d30 remain the delivered outcome. The retained release's prior runtime acceptance is separate from the unpromoted Q6 experiment. This does not mean all ten optimization ideas were implemented or all possible future designs exhausted.

    Restoration passed with no forced cleanup. Shutdown was verified offline at 11:50:57 UTC with fresh 5 W, zero requests and no hold. A later read-only check at 11:58 UTC still reported offline, no loaded models, no requests and no hold. All raw/failed/excluded receipts and local candidate patches are preserved. No idle inhibition, alarm reset or proxy restart was used, and no GPU work remains running or scheduled by this task.

    Remaining work for this bounded task: none. Closing with the accepted historical release and the completed, explicitly limited follow-up results. Future source work can use this evidence without treating unimplemented ideas or the Q6 screen as verified production improvements.
25. novkien

    closed this as completedon Oct 4, 2026
26. novkien

    reopened this on Oct 4, 2026
27. novkien commented on Oct 4, 2026

    

    novkien

    on Oct 4, 2026

    OwnerAuthor

    More actions

    Owner decision and implementation claim - promote the reviewed Q6 candidate

    The owner explicitly instructed Codex in the current task to deploy the Q6 small-batch optimization, accepting a positive gain below the earlier 5% threshold. This supersedes the prior decision not to promote the two-pair screen. The owner has not requested another performance campaign. Reopening this same work issue for delivery of the already reviewed candidate; this is not a new regression.

    Codex /root is accountable for exact pure patch 51b9b224b2ff6d3168b5f6cf91720d5a6018b7c77844bd1cabdb92a018b4f7d5 on base 5423bddf352adc41a8cc59ed02d38a511b333d30, in branch codex/flash9-quant-combined-5423. Scope is the measured SM70 Q6_K W4 shapes plus the measured W1 draft-head shape and their existing backend fixtures. No Q5, profiling overlay, model/GGUF change, route retuning or proxy runtime change is included.

    Reuse the passed native/RPC builds, 100 CPU-reference rows, both actual-model route proofs, and independently validated two-pair screen. Those descriptive gains remain Qwen +3.51%-3.81%; Flash -0.97% to +4.55% with run-to-run variability. Owner acceptance of this candidate does not turn them into five-pair confidence estimates or universal speedup claims.

    Next: root revalidates the unchanged patch and focused local checks; a delivery delegate publishes the reviewed PR and merges the exact head with this open issue linked. Build a new matched native/RPC release from the exact merge, retain 5423 for rollback, activate via normal model lifecycle, and verify source identity plus bounded inference on both models through proxy8082. No repeated TG campaign or unrelated full-context sweep is required for this unchanged dispatch-only candidate. Keep all power constraints: no idle inhibition; wake only for actual build/load/inference, then immediate unload/shutdown with physical verification. Issue stays open until the merged pair is actually installed and the live checks complete.
28. novkien commented on Oct 4, 2026

    

    novkien

    on Oct 4, 2026

    OwnerAuthor

    More actions

    Pre-merge handoff — Q6 candidate

    PR #23: #23

    The PR is based on master at 5423bddf352adc41a8cc59ed02d38a511b333d30, with reviewed head 46871624dfd0b8b28fcc91683b82838d392dedef. The exact base-to-head patch hash is 51b9b224b2ff6d3168b5f6cf91720d5a6018b7c77844bd1cabdb92a018b4f7d5; only ggml/src/ggml-cuda/mmvq.cu and tests/test-backend-ops.cpp changed.

    Root reviewed the full patch and surrounding kernel/security logic; focused style/name/diff checks passed; independent review returned GO on the same patch. Reused evidence: native/RPC builds, 100/100 CPU-reference cases, actual Qwen and Flash route proofs, and the descriptive two-pair-per-model screen. The reported screen remains Qwen +3.51% to +3.81% and Flash -0.97% to +4.55%, without confidence intervals or a universal speedup claim.

    The PR is mergeable. master has no branch protection and the repository has no rulesets; pending/queued CI is not an enforced merge gate. After merge, Codex /root remains responsible for building the exact merge, deploying it, and verifying bounded inference on both models through proxy :8082. This issue remains open until those steps pass.
29. novkien commented on Oct 4, 2026

    

    novkien

    on Oct 4, 2026

    OwnerAuthor

    More actions

    Final handoff - Q6 merged, deployed and live-verified

    The owner's explicit Q6 deployment decision is fulfilled by PR23, exact merge c2d2a20df48b261d4ce94a3b14374ce4bfed9282. Native and RPC packages were built from that merge and are now installed and retained. Root verified running executable/library identities and HTTP200 cold8192+512 inference on BOTH unchanged Qwen27B and FlashNext routes, normal slot release/unload, final installed identities and retained rollback packages. The complete acceptance report separates source/build/install/activation/behavior and records the corrected CLI-only first attempt.

    The earlier two-pair gains remain descriptive; the owner selected them without the former5% promotion threshold. No repeated comparative benchmark or unrelated optimization was added. Final wrapper exit0 confirms accepted Q6 retention and no forced cleanup. Shutdown was physically verified offline at13:10:01UTC on2026-10-04 with fresh6W, no hold and zero requests. All idle-inhibition prohibitions were preserved.

    No remaining work for this selected Q6 delivery. Closing issue9 with the current body updated to the actual deployed revision. The previous no-promotion decision is superseded; unselected and unimplemented optimization families remain honestly scoped in the history.
30. novkien

    closed this as completedon Oct 4, 2026
31. novkien commented on Oct 4, 2026

    

    novkien

    on Oct 4, 2026

    OwnerAuthor

    More actions

    Owner-directed comparison baseline is ready

    The owner asked to prepare the latest stock llama.cpp binary first, preserve the best selected fork release, and let the owner load both backends and compare. The owner also explicitly directed the next native optimization work to implement a candidate, test TG, merge when it improves, and discard neutral/regressing candidates; the former5% promotion threshold is not to be reintroduced. Existing correctness and scope constraints remain. Further experiments wait for this owner-run comparison rather than changing the fork baseline underneath it.

    Stock backend llama.cpp is now built and installed as a matched native/RPC pair at upstream revision 46847e61582097979f539595d893d83d8e1d1af1, with an empty source-diff hash. Native uses CUDA12.9/SM70-real; RPC uses CUDA13.3/SM120a-real. Build, dependency resolution, native version output, installed manifests and file hashes passed. No model was loaded and no inference or comparison was run by the agent; this is binary readiness for the owner's next load, not a claim of model-level acceptance.

    The existing stock checkouts were on restored-20260927 at b23efaa with eleven local source modifications. Those branches, HEADs and edits were preserved. Clean upstream source was fetched into separate worktrees at /home/jarvis/.local/share/llama-builds/upstream-46847e615820/source on both hosts; only the complete installed llama.cpp/build/bin packages were exchanged, with rollback copies retained.

    The fork remains exactly PR23/c2d2a20df48b261d4ce94a3b14374ce4bfed9282 on master and in both installed packages. Full before/after fork manifests and source identity checks matched. Routes, model files and bridge/proxy services were unchanged.

    The completed stock update receipt SHA-256 is cdba18a894497d913746a71a8cb456ca739b84179b066c21610df019e7621d55. The native build finished in the first compile interval; the RPC build continued from its cache in a second bounded interval without idle inhibition. Final shutdown was verified offline at14:49:57UTC on2026-10-04 with fresh3W, no hold and zero requests. The owner's immediate pull/build preparation request is complete. When source optimization resumes after the comparison, recheck the current base and reopen the primary work issue before edits/PRs, reusing the owner's decision above without asking again for the unchanged scope.