From d3abcd2c27c42d5bec4a66c5c139d309e9c5c94c Mon Sep 17 00:00:00 2001 From: bri-prism <288398250+bri-prism@users.noreply.github.com> Date: Fri, 28 Aug 2026 10:37:33 -0700 Subject: [PATCH 1/3] llama: name the pre-rename PQ2_0 ftype instead of reporting unknown (#134) ggufs packed before the Q2_0_G128 -> PQ2_0 rename carry the old ftype value 142. The format did not change, only the enum did, so these files load and compute correctly but print print_info: file type = unknown, may not work which reads as a broken file. With this they print print_info: file type = PQ2_0 - 2.13 bpw (group 128, legacy ftype) Naming only. No load path, no quantizer path, and no behaviour changes; the new constant exists so the switch has a case for the value. Worth noting the collision that makes this confusing: 142 is also GGML_TYPE_PQ2_0, the tensor type id, which is current and correct. Only the ftype was renumbered, to 141. Verified on a published gguf that carries ftype 142: reports unknown on the current branch, reports PQ2_0 with this change, same output either way. --- include/llama.h | 1 + src/llama-model-loader.cpp | 3 +++ 2 files changed, 4 insertions(+) diff --git a/include/llama.h b/include/llama.h index 34b42dc1f458..e424d5c40100 100644 --- a/include/llama.h +++ b/include/llama.h @@ -157,6 +157,7 @@ extern "C" { LLAMA_FTYPE_MOSTLY_Q1_0 = 40, // except 1d tensors LLAMA_FTYPE_MOSTLY_Q2_0 = 41, // except 1d tensors LLAMA_FTYPE_MOSTLY_PQ2_0 = 141, // except 1d tensors (Prism group-128 Q2_0; matches published PQ2_0 ggufs) + LLAMA_FTYPE_MOSTLY_PQ2_0_LEGACY = 142, // pre-rename value for the same format, still found in published ggufs LLAMA_FTYPE_GUESSED = 1024, // not specified in the model file }; diff --git a/src/llama-model-loader.cpp b/src/llama-model-loader.cpp index 7e7eb5cca668..c43499aa234a 100644 --- a/src/llama-model-loader.cpp +++ b/src/llama-model-loader.cpp @@ -40,6 +40,9 @@ const char * llama_ftype_name(llama_ftype ftype) { case LLAMA_FTYPE_MOSTLY_Q1_0: name = LLAMA_FTYPE_PREFIX "Q1_0"; break; case LLAMA_FTYPE_MOSTLY_Q2_0: name = LLAMA_FTYPE_PREFIX "Q2_0"; break; case LLAMA_FTYPE_MOSTLY_PQ2_0: name = LLAMA_FTYPE_PREFIX "PQ2_0 - 2.13 bpw (group 128)"; break; + // ggufs packed before the Q2_0_G128 -> PQ2_0 rename carry the old ftype value. + // They load and compute correctly; name it so it does not report as unknown. + case LLAMA_FTYPE_MOSTLY_PQ2_0_LEGACY: name = LLAMA_FTYPE_PREFIX "PQ2_0 - 2.13 bpw (group 128, legacy ftype)"; break; case LLAMA_FTYPE_MOSTLY_Q4_0: name = LLAMA_FTYPE_PREFIX "Q4_0"; break; case LLAMA_FTYPE_MOSTLY_Q4_1: name = LLAMA_FTYPE_PREFIX "Q4_1"; break; case LLAMA_FTYPE_MOSTLY_Q5_0: name = LLAMA_FTYPE_PREFIX "Q5_0"; break; From d80fd98bc1ee385b149290d9f7ccf22f9244f43d Mon Sep 17 00:00:00 2001 From: bri-prism <288398250+bri-prism@users.noreply.github.com> Date: Mon, 31 Aug 2026 15:24:33 -0700 Subject: [PATCH 2/3] hadamard: share the activation transform between folded weights that share an input (#137) * hadamard: share the activation transform between folded weights that share an input Folded weights that consume the same activation were each building their own identical transform. ggml does no common-subexpression elimination, so the sign-multiply and the FWHT ran once per weight rather than once per activation. Three groups share an input on this architecture: ffn_gate + ffn_up build_ffn LLM_FFN_PAR, both take cur 64 redundant attn_qkv + attn_gate qwen35.cpp:237,241, both take input 48 redundant attn_q + attn_k + attn_v qwen35.cpp:270,282,285, all take cur 32 redundant That is 144 of 401 folded weights doing a transform that already exists, so 36% of the FWHT calls per token were duplicates. build_lora_mm and build_lora_mm_id now memoize on (input, rotation). Both must match for the transform to be shared: same input tensor and same rotation, which also implies the same width and therefore the same sign vector. The map lives for one graph build. Measured on M5, Metal, a folded low-bit checkpoint, three interleaved pairs on a warm machine: pp512 358.29 -> 370.74 +3.5% tg128 23.20 -> 23.91 +3.1% Run-to-run spread was at most 1.1 on pp512 and 0.11 on tg128, so the gap is well outside the noise. Greedy output is unchanged, byte for byte, on the same prompt. The size of the win matches the prediction: the transform stack was about 10% of traced GPU time, and removing 36% of it should be worth about 3.5%. * hadamard: trim the memo comment to two lines per AGENTS.md --- src/llama-graph.cpp | 16 ++++++++++++++++ src/llama-graph.h | 4 ++++ 2 files changed, 20 insertions(+) diff --git a/src/llama-graph.cpp b/src/llama-graph.cpp index 49c865e7fcf4..f7c14f573414 100644 --- a/src/llama-graph.cpp +++ b/src/llama-graph.cpp @@ -1512,6 +1512,12 @@ ggml_tensor * llm_graph_context::build_lora_mm( const auto it = hadamard_rotations->find(w); if (it != hadamard_rotations->end()) { const auto & t = it->second; + // another folded weight on this same activation already built the transform + const auto memo_key = std::make_pair((const ggml_tensor *) cur, (const ggml_tensor *) t.rot); + const auto memo_it = hadamard_memo.find(memo_key); + if (memo_it != hadamard_memo.end()) { + cur_mm = memo_it->second; + } else { if (t.perm_rep > 1) { // tiled [hd, nk, rep] -> grouped [hd, rep, nk] feature order ggml_tensor * x = ggml_is_contiguous(cur_mm) ? cur_mm : ggml_cont(ctx0, cur_mm); @@ -1524,6 +1530,8 @@ ggml_tensor * llm_graph_context::build_lora_mm( cur_mm = ggml_mul(ctx0, cur_mm, t.signs); } cur_mm = llama_mul_mat_hadamard(ctx0, cur_mm, t.rot); + hadamard_memo[memo_key] = cur_mm; + } } } @@ -1564,6 +1572,12 @@ ggml_tensor * llm_graph_context::build_lora_mm_id( const auto it = hadamard_rotations->find(w); if (it != hadamard_rotations->end()) { const auto & t = it->second; + // another folded weight on this same activation already built the transform + const auto memo_key = std::make_pair((const ggml_tensor *) cur, (const ggml_tensor *) t.rot); + const auto memo_it = hadamard_memo.find(memo_key); + if (memo_it != hadamard_memo.end()) { + cur_mm = memo_it->second; + } else { if (t.perm_rep > 1) { // tiled [hd, nk, rep] -> grouped [hd, rep, nk] feature order ggml_tensor * x = ggml_is_contiguous(cur_mm) ? cur_mm : ggml_cont(ctx0, cur_mm); @@ -1576,6 +1590,8 @@ ggml_tensor * llm_graph_context::build_lora_mm_id( cur_mm = ggml_mul(ctx0, cur_mm, t.signs); } cur_mm = llama_mul_mat_hadamard(ctx0, cur_mm, t.rot); + hadamard_memo[memo_key] = cur_mm; + } } } diff --git a/src/llama-graph.h b/src/llama-graph.h index 492574f69a18..de4674b70b7d 100644 --- a/src/llama-graph.h +++ b/src/llama-graph.h @@ -1078,6 +1078,10 @@ struct llm_graph_context { const llama_hadamard_rotations * hadamard_rotations; const llama_hadamard_rotations * hadamard_inverses; + // Transforms shared by folded weights on the same activation. Key is (input, rotation); + // both must match. Valid for one graph build only. + mutable std::map, ggml_tensor *> hadamard_memo; + std::map samplers; const llm_graph_cb & cb_func; From a2c7004f54ac73d8769b1d5fbe251bbfd8f0b58c Mon Sep 17 00:00:00 2001 From: bri-prism <288398250+bri-prism@users.noreply.github.com> Date: Fri, 4 Sep 2026 12:56:48 -0700 Subject: [PATCH 3/3] ggml: add PTQ1_0, ternary at group 128 (1.75 bpw, lossless vs PQ2_0) (#148) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * ggml: add PTQ1_0, Prism ternary at group 128 (1.75 bpw) Upstream TQ1_0 packs ternary weights base-3 at 5 trits per byte, but carries a single fp16 scale per 256 weights. The ternary checkpoints this targets carry a scale per 128, so a 256-wide scale has to discard one of the two group scales it straddles. Measured on a ternary checkpoint of this family: every trit survives and the zero pattern is preserved exactly, but the smaller-alpha group of each block is reconstructed at the larger alpha, giving 6.43% weight error that passes into activations at full strength because it is multiplicative on the weight rather than noise. PTQ1_0 keeps TQ1_0's trit packing and our group of 128: qs[24] + qh[2] + fp16 d = 28 bytes / 128 weights = 1.7500 bpw against PQ2_0's 34 bytes = 2.1250 bpw, so 17.6% smaller with no loss of information: PQ2_0 spends two bits on a value carrying log2(3) = 1.585 bits. test-quantize-fns reports ptq1_0 identical to pq2_0 on both metrics, absolute quantization error 0.008678 and dot product error 0.141111, versus tq1_0's 0.008681 and 0.141345. A round trip of real folded ternary weights is bit-identical to the PQ2_0 reconstruction, so the only error either format carries is the shared fp16 scale rounding. A distinct type id (143) rather than changing TQ1_0 in place, for the same reason PQ2_0 is distinct from upstream's group-64 Q2_0: TQ1_0 is a wire format, and reinterpreting its block width would silently misread every existing TQ1_0 gguf. TQ1_0's qs staging is fixed at 32-then-16 bytes, which cannot cover 24, so the stages are generalised to 32/16/8. At TQ1_0's 48-byte qs that reduces to exactly its original behaviour. Scope: codec, ggml traits, CPU vec_dot against Q8_0, and the ftype plumbing so llama-quantize can target it. The vec_dot decodes a block to element order before the dot, sharing the traversal with dequantize so the two cannot drift; it is generic C on every arch, with no SIMD path yet. CUDA, Metal and Vulkan kernels, and the repack gemv/gemm paths, are not implemented. * ggml: wire PTQ1_0 into quantize dispatch, ftype map and row validation Three plumbing gaps the traits entry alone did not cover, each of which failed end-to-end quantization in a different way: ggml_quantize_chunk - without a dispatch case the returned byte count did not match the row size and GGML_ASSERT(result == nrows * row_size) aborted a worker thread mid-tensor, leaving a truncated output file behind. ggml_ftype -> wtype - the ftype could not resolve to a ggml type. validate_row_data - "invalid type 143", so quantization refused the result it had just produced. Verified end to end on a 1.7B ternary model with the embedding and output tensors pinned to f16 in both runs, so only the quantized tensors differ: quantized tensors PQ2_0 357.47 MiB -> PTQ1_0 294.47 MiB, 17.62% smaller (1.75/2.125 bpw predicts 17.65%) perplexity, 4 chunks at ctx 512, identical to every printed digit: per chunk 11.3796 / 14.7182 / 14.5014 / 12.6944 for both, final 12.6944 +/- 1.12596 for both The perplexity run goes through the new CPU vec_dot, so that path is exercised rather than only the codec. * metal: PTQ1_0 kernels — correct, but 4.2x slower decode than PQ2_0 Adds the Metal path for PTQ1_0: an element accessor for the base-3 layout, both dequantize entry points, a mat-vec kernel, mul_mm and mul_mm_id instantiations for f32 and f16, get_rows, the simdgroup constants and both pipeline-selection switches. Also enumerates PTQ1_0 in test-backend-ops, which otherwise silently skipped the new type and reported a clean pass. Correct: 1251/1251 MUL_MAT tests pass on Apple M5 Pro. End to end on a 1.7B ternary model the Metal perplexity is identical to PQ2_0 to every printed digit, 12.6935 +/- 1.12567 for both. Slow, and the reason is structural. Measured on M5 Pro: PQ2_0 PTQ1_0 pp128 4505 tok/s 2590 tok/s 0.57x tg32 182.5 tok/s 43.6 tok/s 0.24x Base-3 packing makes element order non-positional, so a thread owning 16 contiguous weights cannot index bytes directly, and threads owning different trit indices reload the same bytes. Hoisting the trit factor out of the inner loop for spans below element 80 took decode from 29.8 to 43.6 tok/s, but the remaining gap needs the thread mapping changed so one thread consumes all five trits of a byte it loads once, instead of five threads each loading it for one trit. This reproduces what the earlier TQ1_0 work already found on other backends: base-3 is ALU-bound rather than bandwidth-bound, and the packing win does not pay for itself at decode without that restructure. The 17.6% file-size reduction is real and lossless; the speed is not there yet. * metal: PTQ1_0 decode 29.8 -> 165.3 tok/s, now 0.89x of PQ2_0 Four changes, each measured on Apple M5 Pro against a 1.7B ternary model, with MUL_MAT staying 1251/1251 and Metal perplexity identical to PQ2_0 throughout: 1. byte-owning threads 29.8 -> 136.4 a thread now consumes all five trits of a byte it loads once, instead of eight threads each reloading the same qs bytes for one trit index; the block's 26 bytes are read once total 2. y staged in registers 136.4 -> 143.5 y was being re-read from device memory once per row; it is now gathered once per block and reused, with the ternary offset folded into a single sumy subtraction 3. N_R0 8 -> 4 143.5 -> 163.8 rows-per-thread was the real limiter: 16 rows gives 94 tok/s, 8 gives 144, 4 gives 164, 2 gives 163 4. balanced qh 163.8 -> 165.3 qh's 8 elements went to threads 0 and 1 only, so those did 19 elements while the rest did 15 and the whole simdgroup waited; one element per thread makes it 16 each and drops yl to 16 Trit extraction is the base-3 remainder recurrence, t = (v*3)>>8 with v = (v*3)&0xFF, two integer ops per trit with no table. A 256-entry lookup was tried and is not better for decode: it adds a dependent load per byte, 26 extra per block against PQ2_0's 32 total loads. It does help the mul_mm path, which must produce 16 contiguous elements and cannot use the byte-owning mapping, so dequantize keeps the table. Final, PQ2_0 vs PTQ1_0 on M5 Pro: pp128 4541 -> 3998 tok/s 0.88x tg32 186 -> 165 tok/s 0.89x size 357 -> 294 MiB 17.6% smaller, and lossless Prefill's remaining gap is structural: mul_mm's dequantize contract is 16 contiguous elements, so five threads still touch the same 16 bytes for different trit indices. Closing it needs mul_mm itself to stage a decoded block in threadgroup memory, which is a larger change than this. * cuda: PTQ1_0 dequantize and MMVQ paths — NOT COMPILED, no toolchain here Adds the CUDA surface for PTQ1_0, mirroring PQ2_0: type traits, a device trit accessor, the float2 dequantize entry point, the six dequant dispatch sites in convert.cu, get_rows, the MMVQ dot with its VDR constant and three dispatch sites, and both supports_op type lists. **This code has never been compiled or executed.** The machine it was written on is an Apple M5 Pro with no CUDA toolchain, so nvcc was unavailable. It must be built and run on a CUDA host before anyone relies on it. What *is* verified is the part where a bug would actually hide. The device accessor's element mapping is pure integer logic, so it was transcribed verbatim to the host and checked against the CPU codec's own traversal over 2,560,000 element positions across 20,000 random blocks: zero mismatches. That check is kept as tests/test-ptq1_0-element-map.cpp so it can be rerun. Trit extraction uses the same base-3 remainder recurrence as the Metal path, t = (v*3)>>8 with v = (v*3)&0xFF. The MMVQ dot gathers four trits at a time and packs them into an int as signed int8 lanes so accumulation still goes through dp4a, since base-3 packing leaves no byte-aligned run inside a 32-element chunk. MMQ is deliberately not implemented. Types without it fall back to dequantize plus cuBLAS for prefill, which is correct if not optimal, and writing an MMQ tile loader blind against tile layouts I cannot test would be worse than leaving the fallback in place. * metal: tune PTQ1_0 to N_SG=1, and record what optimising it ruled out N_R0=4, N_SG=1. Decode lands at 162-166 tok/s across runs against PQ2_0's 186, so 0.87-0.89x; the spread between N_SG 1 and 2 is inside run-to-run noise (+/- 3.4 on a single measurement). Five further optimisations were tried and MEASURED, and four of them lost. Kept here so nobody spends the same day twice: 256-entry trit lookup table 142.1 vs 143.5 recurrence no decode gain (helps mul_mm, which keeps it) independent pow3 extraction 159.8 vs 166.0 recurrence WORSE contiguous byte ownership 162.0 vs 165.3 scattered WORSE N_R0 = 3 / 5 / 6 129.8 / 156.7 / 149.6 all worse than 4 N_SG = 8 161.1 worse than 1-2 What that rules out. The lookup table shows decode is not bound by trit arithmetic. The pow3 variant shows it is not bound by the recurrence's five-deep serial dependency either -- four rows per thread already supply enough independent work, so fewer ops beat more ILP. Contiguous byte runs made it worse because thread 5 straddles the qs 16/8 chunk boundary, and the resulting per-byte branch cost more than consolidating the loads saved. Where the remaining 11% actually lives, per 128 weights: PQ2_0 128 x (shift, mask, cvt, fma) = 512 ops, 34 B read PTQ1_0 128 x (mul, shift, cvt, fma, mask-for-next) = 640 ops, 28 B read ALU 1.25x, bytes 0.82x A purely ALU-bound kernel would sit at 0.80x and a purely bandwidth-bound one at 1.21x. Measured 0.89x sits between them, which is about where this instruction and byte mix should land. Extracting a base-3 digit is inherently a multiply where a 2-bit field is a shift, so this gap is the price of packing at the 1.585-bit information limit rather than on a 2-bit boundary, not an implementation defect. The one lever left with a real mechanism behind it is prefill, not decode. In decode each row has its own weights so there is nothing to share, but mul_mm reuses a weight tile across many output columns, and its dequantize contract forces 16 contiguous elements, so five threads still touch the same 16 bytes. Staging a cooperatively decoded block in threadgroup memory would remove that redundancy outright. Prefill is at 0.87x (3946 vs 4492). Perplexity remains identical to PQ2_0 at 12.6935 +/- 1.12567, MUL_MAT 1251/1251. * metal, cpu: close PTQ1_0 type-list gaps, add the missing mul_mv_ext r1_5 Self-review compared every switch that names PQ2_0 against those naming PTQ1_0 and found the new type absent from ten support enumerations: seven in ggml-cpu/ops.cpp, including ggml_compute_forward_get_rows_f32, one in ggml-metal-ops.cpp and two in ggml-metal-device.m. PTQ1_0 is traits-compatible with PQ2_0 there -- same block size, same to_float and from_float shape -- so it belongs in each. Closing those gaps then exposed a real defect. Adding the type to the Metal support lists enabled the mul_mv_ext path for wider n, which immediately failed with "kernel not found in any metal library" for kernel_mul_mv_ext_ptq1_0_f32_r1_5: PQ2_0 instantiates r1_2 through r1_5 and only r1_2 through r1_4 had been written, because the grep that enumerated them was truncated by head. Instantiating r1_5 fixes it, and the instantiation set is now derived from PQ2_0's rather than transcribed. Worth noting the earlier 1251/1251 passes were not wrong, they were narrow: with the type missing from the support lists the wider-n path was never selected, so the absent kernel could not be reached. The gap and the defect concealed each other. After: MUL_MAT 1251/1251 on M5 Pro, Backend MTL0 OK, Metal perplexity identical to PQ2_0 at 12.6935 +/- 1.12567, decode 165 tok/s and prefill 3947. * tests: verify the CUDA MMVQ dot numerically on the host The CUDA path still cannot be compiled here, but its arithmetic can be checked. This stubs the three intrinsics the kernel uses -- __dp4a, __low2float and the 4-byte int load -- transcribes vec_dot_ptq1_0_q8_1 and the device trit accessor verbatim, and compares them against a float reference built by dequantising the block and the q8_1 activations independently. Result over 5000 random weight blocks with random q8_1 activations: exact integer-accumulator checks 20000, mismatches 0 error / sum of term magnitudes 2.399e-08 The integer test is the one that matters: the kernel accumulates in int via dp4a before scaling, so if the trit extraction, the int8 lane packing or the chunk-to- q8_1 pairing were wrong, that accumulator would differ. It does not, on any of the 20000 chunks. A first pass reported 7.062e-05 "relative error" and looked like a failure. That was cancellation: the reference total is a signed sum that can land near zero, so dividing by it inflates the ratio. Measured against the sum of term magnitudes instead, which cancellation cannot distort, the error is 2.4e-08, i.e. float32 rounding of the four per-chunk products. Together with tests/test-ptq1_0-element-map.cpp this leaves only compilation unproven for the CUDA path, not correctness of its math. * cuda: move ptq1_0_trit to common.cuh so vecdotq.cuh can see it — BUILDS NOW First compile of the CUDA path, on an RTX 5090 with CUDA 12.8 and CMAKE_CUDA_ARCHITECTURES=120, failed with vecdotq.cuh(822): error: identifier "ptq1_0_trit" is undefined The accessor was defined in dequantize.cuh, which vecdotq.cuh does not include. PQ2_0's dot never needed a shared helper -- it decodes 2-bit fields inline with byte_perm -- so nothing in the file it was mirrored from exposed the dependency. Both headers include common.cuh, so the helper belongs there. Verified on the 5090 after the fix: test-backend-ops -o MUL_MAT 1283/1283 passed, Backend CUDA0: OK 45 ptq1_0 cases OK, 0 failures host element-map test 2,560,000 positions, 0 mismatches host MMVQ dot test 20,000 integer checks, 0 mismatches The two host tests were already green before this, which is the point worth recording: they verified the arithmetic and could not have caught a missing declaration. Compilation was the only gate that would find this, and the earlier commit saying the CUDA path had never been built was accurate about the risk. * vulkan: PTQ1_0 support, mirrored from Q1_0 not Q2_0 — UNTESTED as committed Adds the Vulkan path for PTQ1_0: the block struct and QUANT_K defines, the trit accessor plus dequantize/dequantize4/get_dm, the mul_mm load branch, a standalone dequant shader, the shader-generator registration, and 16 pipeline sites in ggml-vulkan.cpp mirrored from Q1_0's. Mirrored from Q1_0 deliberately. Q1_0 is type 41 at group 128 with correct Vulkan kernels; Q2_0 in this tree is group 128 in ggml-common.h but its Vulkan side still declares QUANT_K_Q2_0 64, because upstream's group-64 kernels were merged over a group-128 type and only some of the GLSL was re-derived. Pattern-matching off Q2_0 would have inherited that. Q1_0 also has no MMVQ entry, so this follows the same dequantize-based path rather than inventing one. One divergence from Q1_0 worth naming: PTQ1_0 puts the scale LAST in the block (qs[24], qh[2], then d) where Q1_0 puts d first. The GLSL struct mirrors the C field order exactly. Getting that wrong misindexes every block silently instead of failing, which is the same class of error as the Q2_0 breakage above. Not wired, and it is a real gap rather than an oversight: on-device quantization (cpy_f32_quant, set_rows, cpy_quant_f32). Those need a base-3 ENCODER in GLSL, which is a separate piece of work and is not needed to run a model -- weights are quantized offline by llama-quantize. **This has not been compiled or run.** There was no Vulkan toolchain on the machine it was written on. It is pushed so it can be built and tested on a GPU box, and it should be treated as unverified until test-backend-ops says otherwise. * cuda: PTQ1_0 MMVQ takes the whole block — measured, and still 6-9x slower Benchmarked the CUDA path for the first time. It is correct and slow, and the numbers matter more than the change: us/run on a 4090 at m=4096, k=14336, against pq2_0 on the same run. n=1 n=2 n=3 n=8 n=512 q1_0 8.88 9.82 13.05 30.14 225.4 pq2_0 9.75 11.05 14.92 31.27 232.4 ptq1_0 46.8 64.6 99.1 287.3 525.6 The per-chunk version asked "which byte holds element e" per element, a branch plus a variable-length recurrence, and iqs differs per lane so the branch diverged. Taking the whole 128-weight block per call (VDR 1 -> 4) makes every lane walk the same 26 bytes in the same order. That bought about 9% at n=2, 70.6 -> 64.6, and no more. So divergence was not the dominant cost, and the comment now says so. What remains is load shape: this issues 128 scalar int8 reads of the q8_1 activations, one per element, where pq2_0 issues eight dp4a calls over 4-byte words. The fix is identified but deliberately not guessed at. Below element 80 the element stride is 16 and 4 divides 16, so four consecutive elements are four consecutive BYTES at the same trit index. Four recurrences packed into one int would feed dp4a against a 4-byte q8_1 load, matching pq2_0's shape. That needs a build cycle to verify, and this session has already shown twice -- on the CPU vec_dot and on the Metal mapping -- that predictions about where GPU time goes are wrong about as often as they are right. Correctness held throughout: test-backend-ops -o MUL_MAT 1283/1283 on both an RTX 5090 and an RTX 4090, Backend CUDA0 OK, 45 ptq1_0 cases OK and 0 bad, before and after. Note n=512 is the cuBLAS dequantize path rather than MMVQ, and is 2.3x rather than 6-9x, which is consistent with the problem being in the MMVQ load pattern specifically. * vulkan: fix three build breaks — shared ptq1_0.glsl, skip coopmat2, drop cm2 pipelines The Vulkan commit did not build. Three errors, one root cause each, all found by a teammate on an Intel Arc B390 and independently by windows CI: 1. mul_mm_cm2.comp: 'dequantFuncA' undeclared. Adding the type to type_names in vulkan-shaders-gen.cpp auto-enrols it in coopmat2 generation, but dequant_funcs_cm2.glsl has no PTQ1_0 decoder (Q1_0 has six entries). That failed the whole Vulkan shader build, not just this type. Now skipped for ptq1_0, which falls back to the scalar and coopmat1 paths that are actually implemented. 2. mul_mm_funcs.glsl: 'ptq1_0_trit' no matching function. The helper lived in dequant_funcs.glsl; mul_mm.comp includes types/dot_product/mul_mm_id/mul_mm_funcs and never dequant_funcs.glsl. **This is the same bug as the CUDA one fixed an hour earlier, same helper, different backend**: defined in the dequantize chain, called from the matmul chain. Fixed the same way, with a shared header vulkan-shaders/ptq1_0.glsl behind an include guard, included from both. Not duplicated: two copies would have to stay in step with the CPU codec in ggml-quants.c, and a divergence there appears as wrong matmul results, never as a build error. 3. ggml-vulkan.cpp: 'matmul_ptq1_0_f16_cm2_len' undeclared. Registrations for cm2 pipelines whose shaders no longer exist once (1) is fixed. Both removed, with a comment recording that the absence is deliberate. (1) and (3) had to be fixed together or one error just trades for the other. Verified locally this time, which is the other half of the lesson: glslc, molten-vk, vulkan-headers and vulkan-loader were already installed on the machine that wrote the original commit. It was pushed unverified on the false assumption that no toolchain was present, and a local build finds all three in minutes. -DGGML_VULKAN=ON builds clean SPIR-V generated 1991 shaders, 54 of them ptq1_0, incl. coopmat1, no cm2 test-backend-ops MUL_MAT 1009/1009, Backend Vulkan0: OK, no regression That local run does NOT verify the kernel. MoltenVK reports every ptq1_0 case "not supported", and it reports pq2_0 the same way (q1_0 gets 28 of 156), so this is an Apple/MoltenVK coverage limit rather than a defect in the new type. Kernel correctness comes from the Arc B390 run on the patched tree: 1009/1009 with 78 ptq1_0 cases actually executed and 0 failures, on KHR_coopmat. Still absent by design: MMVQ (no q8_1 pipelines, absent from should_use_mmvq), coopmat2, and on-device quantization, which needs a base-3 encoder in GLSL. * vulkan: make PTQ1_0 reachable, and guard the coopmat2 path The PTQ1_0 pipelines were created but the type was absent from every selection and support enumeration, so supports_op returned false, the scheduler never offered MUL_MAT/MUL_MAT_ID/GET_ROWS to Vulkan, and the generated shaders were unreachable. The build looked green while nothing executed on the backend. Add the type alongside GGML_TYPE_Q1_0 in the five pipeline getters and the two supports_op gates. Adding it there alone is not safe, because PTQ1_0 is the only type in those lists with no coopmat2 pipelines: dequant_funcs_cm2.glsl carries no decoder for it, so vulkan-shaders-gen skips cm2 generation. The cm2 slots are still allocated by vk_matmul_pipeline2's constructor, so the mm getter would return a non-null pipeline whose l/m/s are all null, and the caller's "mmp == nullptr" fallback would not fire; the mul_mat_id getter would instead trip GGML_ASSERT(support_fp32acc). Both getters now return nullptr when the cm2 slots are empty, which routes the caller to the dequant + f16 matmul fallback it already implements. That path needs ggml_vk_get_to_fp16, which is one of the gates opened above. The cpy, SET_ROWS and DUP gates are deliberately left alone: copy_to_quant.comp has no PTQ1_0 block, so quantizing to this type on device is unimplemented and a refusal there is correct. Verified on Intel Arc B390 (KHR_coopmat): ptq1_0 goes from 0 OK / 78 declined to 28 OK / 50 declined / 0 failures, matching q1_0 and q2_0 case for case, byte-identical across two runs. The coopmat2 guard is reasoned from the pipeline-creation structure and is not covered by that run, which is a coopmat1 device; it has no coopmat2 hardware behind it yet. * cuda: accelerate PTQ1_0 MMVQ with packed DP4A Assisted-by: OpenAI Codex --- ggml/include/ggml.h | 4 +- ggml/src/ggml-common.h | 15 ++ ggml/src/ggml-cpu/arch-fallback.h | 15 ++ ggml/src/ggml-cpu/ggml-cpu.c | 6 + ggml/src/ggml-cpu/ops.cpp | 7 + ggml/src/ggml-cpu/quants.c | 72 ++++++++ ggml/src/ggml-cpu/quants.h | 3 + ggml/src/ggml-cuda/common.cuh | 34 ++++ ggml/src/ggml-cuda/convert.cu | 12 ++ ggml/src/ggml-cuda/dequantize.cuh | 10 ++ ggml/src/ggml-cuda/getrows.cu | 4 + ggml/src/ggml-cuda/ggml-cuda.cu | 2 + ggml/src/ggml-cuda/mmvq.cu | 54 +++++- ggml/src/ggml-cuda/vecdotq.cuh | 154 ++++++++++++++++- ggml/src/ggml-metal/ggml-metal-device.cpp | 10 ++ ggml/src/ggml-metal/ggml-metal-device.m | 2 + ggml/src/ggml-metal/ggml-metal-impl.h | 2 + ggml/src/ggml-metal/ggml-metal-ops.cpp | 1 + ggml/src/ggml-metal/kernels/dequantize.h | 83 +++++++++ ggml/src/ggml-metal/kernels/mul_mm.metal | 4 + ggml/src/ggml-metal/kernels/mul_mv.metal | 158 ++++++++++++++++++ ggml/src/ggml-metal/kernels/quantize.metal | 1 + ggml/src/ggml-quants.c | 105 ++++++++++++ ggml/src/ggml-quants.h | 3 + ggml/src/ggml-vulkan/ggml-vulkan.cpp | 41 ++++- .../vulkan-shaders/dequant_funcs.glsl | 21 +++ .../vulkan-shaders/dequant_ptq1_0.comp | 54 ++++++ .../ggml-vulkan/vulkan-shaders/mul_mm.comp | 3 + .../vulkan-shaders/mul_mm_funcs.glsl | 15 ++ .../ggml-vulkan/vulkan-shaders/ptq1_0.glsl | 38 +++++ .../src/ggml-vulkan/vulkan-shaders/types.glsl | 21 +++ .../vulkan-shaders/vulkan-shaders-gen.cpp | 10 +- ggml/src/ggml.c | 10 ++ gguf-py/gguf/constants.py | 9 +- include/llama.h | 1 + src/llama-model-loader.cpp | 2 + src/llama-quant.cpp | 4 +- tests/test-backend-ops.cpp | 6 +- tests/test-ptq1_0-cuda-dot.cpp | 132 +++++++++++++++ tests/test-ptq1_0-element-map.cpp | 64 +++++++ tests/test-quantize-fns.cpp | 3 +- tools/quantize/quantize.cpp | 1 + 42 files changed, 1174 insertions(+), 22 deletions(-) create mode 100644 ggml/src/ggml-vulkan/vulkan-shaders/dequant_ptq1_0.comp create mode 100644 ggml/src/ggml-vulkan/vulkan-shaders/ptq1_0.glsl create mode 100644 tests/test-ptq1_0-cuda-dot.cpp create mode 100644 tests/test-ptq1_0-element-map.cpp diff --git a/ggml/include/ggml.h b/ggml/include/ggml.h index 5c04f2c81874..8fa52d812754 100644 --- a/ggml/include/ggml.h +++ b/ggml/include/ggml.h @@ -433,7 +433,8 @@ extern "C" { // Prism-private Q2_0 at group size 128 (upstream Q2_0 is group 64). High id so it // slots above upstream types; type_traits is sized to COUNT (143) with 43..141 unused. GGML_TYPE_PQ2_0 = 142, - GGML_TYPE_COUNT = 143, + GGML_TYPE_PTQ1_0 = 143, // Prism-private ternary, group 128 + GGML_TYPE_COUNT = 144, }; // precision @@ -479,6 +480,7 @@ extern "C" { GGML_FTYPE_MOSTLY_Q1_0 = 27, // except 1d tensors GGML_FTYPE_MOSTLY_Q2_0 = 28, // except 1d tensors GGML_FTYPE_MOSTLY_PQ2_0 = 128, // except 1d tensors (Prism-private group-128 Q2_0) + GGML_FTYPE_MOSTLY_PTQ1_0 = 129, // except 1d tensors (Prism-private group-128 ternary) }; // available tensor operations: diff --git a/ggml/src/ggml-common.h b/ggml/src/ggml-common.h index 7cde4dd68bf2..65075cad860c 100644 --- a/ggml/src/ggml-common.h +++ b/ggml/src/ggml-common.h @@ -101,6 +101,8 @@ typedef sycl::half2 ggml_half2; #define QI_PQ2_0 (QK_PQ2_0 / 32) #define QR_PQ2_0 1 +#define QI_PTQ1_0 (QK_PTQ1_0 / 32) +#define QR_PTQ1_0 1 #define QI4_0 (QK4_0 / (4 * QR4_0)) @@ -204,6 +206,19 @@ typedef struct { } block_pq2_0; static_assert(sizeof(block_pq2_0) == sizeof(ggml_half) + QK_PQ2_0 / 4, "wrong pq2_0 block size/padding"); +// PTQ1_0: Prism-private ternary at group size 128. Same base-3 trit packing as +// upstream TQ1_0 (type 34) but one fp16 scale per 128 weights instead of per 256. +// 1.75 bpw vs PQ2_0's 2.125, and lossless for checkpoints that are already ternary +// at group 128 -- TQ1_0 cannot represent those, because a 256-wide scale has to +// discard one of the two group scales it straddles. +#define QK_PTQ1_0 128 +typedef struct { + uint8_t qs[(QK_PTQ1_0 - 4*QK_PTQ1_0/64)/5]; // 24 B, 5 trits per byte -> 120 values + uint8_t qh[QK_PTQ1_0/64]; // 2 B, 4 trits per byte -> 8 values + ggml_half d; // scale +} block_ptq1_0; +static_assert(sizeof(block_ptq1_0) == sizeof(ggml_half) + QK_PTQ1_0/64 + (QK_PTQ1_0 - 4*QK_PTQ1_0/64)/5, "wrong ptq1_0 block size/padding"); + #define QK4_0 32 typedef struct { ggml_half d; // delta diff --git a/ggml/src/ggml-cpu/arch-fallback.h b/ggml/src/ggml-cpu/arch-fallback.h index f9f4930fa452..d283c2ed51be 100644 --- a/ggml/src/ggml-cpu/arch-fallback.h +++ b/ggml/src/ggml-cpu/arch-fallback.h @@ -18,6 +18,7 @@ #define ggml_vec_dot_nvfp4_q8_0_generic ggml_vec_dot_nvfp4_q8_0 #define ggml_vec_dot_q1_0_q8_0_generic ggml_vec_dot_q1_0_q8_0 #define ggml_vec_dot_pq2_0_q8_0_generic ggml_vec_dot_pq2_0_q8_0 +#define ggml_vec_dot_ptq1_0_q8_0_generic ggml_vec_dot_ptq1_0_q8_0 #define ggml_vec_dot_q2_0_q8_0_generic ggml_vec_dot_q2_0_q8_0 #define ggml_vec_dot_tq1_0_q8_K_generic ggml_vec_dot_tq1_0_q8_K #define ggml_vec_dot_tq2_0_q8_K_generic ggml_vec_dot_tq2_0_q8_K @@ -79,6 +80,8 @@ #define ggml_gemm_q1_0_4x8_q8_0_generic ggml_gemm_q1_0_4x8_q8_0 #define ggml_gemm_pq2_0_4x8_q8_0_generic ggml_gemm_pq2_0_4x8_q8_0 #elif defined(__aarch64__) || defined(__arm__) || defined(_M_ARM) || defined(_M_ARM64) +// PTQ1_0 currently has only the generic vec_dot; alias it here until a SIMD version lands +#define ggml_vec_dot_ptq1_0_q8_0_generic ggml_vec_dot_ptq1_0_q8_0 // repack.cpp #define ggml_quantize_mat_q8_K_4x4_generic ggml_quantize_mat_q8_K_4x4 #define ggml_quantize_mat_q8_K_4x8_generic ggml_quantize_mat_q8_K_4x8 @@ -91,6 +94,8 @@ #define ggml_gemm_q2_K_8x8_q8_K_generic ggml_gemm_q2_K_8x8_q8_K #define ggml_gemm_pq2_0_4x8_q8_0_generic ggml_gemm_pq2_0_4x8_q8_0 #elif defined(__x86_64__) || defined(__i386__) || defined(_M_IX86) || defined(_M_X64) +// PTQ1_0 currently has only the generic vec_dot; alias it here until a SIMD version lands +#define ggml_vec_dot_ptq1_0_q8_0_generic ggml_vec_dot_ptq1_0_q8_0 // quants.c #define ggml_vec_dot_q2_0_q8_0_generic ggml_vec_dot_q2_0_q8_0 // repack.cpp @@ -121,6 +126,8 @@ #define ggml_gemm_q8_0_4x8_q8_0_generic ggml_gemm_q8_0_4x8_q8_0 #define ggml_gemm_q1_0_4x4_q8_0_generic ggml_gemm_q1_0_4x4_q8_0 #elif defined(__POWERPC__) || defined(__powerpc__) +// PTQ1_0 currently has only the generic vec_dot; alias it here until a SIMD version lands +#define ggml_vec_dot_ptq1_0_q8_0_generic ggml_vec_dot_ptq1_0_q8_0 // ref: https://github.com/ggml-org/llama.cpp/pull/14146#issuecomment-2972561679 // quants.c #define quantize_row_q8_K_generic quantize_row_q8_K @@ -175,6 +182,8 @@ #define ggml_gemm_q1_0_4x8_q8_0_generic ggml_gemm_q1_0_4x8_q8_0 #define ggml_gemm_pq2_0_4x8_q8_0_generic ggml_gemm_pq2_0_4x8_q8_0 #elif defined(__loongarch64) +// PTQ1_0 currently has only the generic vec_dot; alias it here until a SIMD version lands +#define ggml_vec_dot_ptq1_0_q8_0_generic ggml_vec_dot_ptq1_0_q8_0 // quants.c #define quantize_row_q8_K_generic quantize_row_q8_K #define ggml_vec_dot_tq1_0_q8_K_generic ggml_vec_dot_tq1_0_q8_K @@ -229,6 +238,8 @@ #define ggml_gemm_q1_0_4x8_q8_0_generic ggml_gemm_q1_0_4x8_q8_0 #define ggml_gemm_pq2_0_4x8_q8_0_generic ggml_gemm_pq2_0_4x8_q8_0 #elif defined(__riscv) +// PTQ1_0 currently has only the generic vec_dot; alias it here until a SIMD version lands +#define ggml_vec_dot_ptq1_0_q8_0_generic ggml_vec_dot_ptq1_0_q8_0 // quants.c #define ggml_vec_dot_nvfp4_q8_0_generic ggml_vec_dot_nvfp4_q8_0 #define ggml_vec_dot_q2_0_q8_0_generic ggml_vec_dot_q2_0_q8_0 @@ -276,6 +287,8 @@ #define ggml_gemm_q1_0_4x8_q8_0_generic ggml_gemm_q1_0_4x8_q8_0 #define ggml_gemm_pq2_0_4x8_q8_0_generic ggml_gemm_pq2_0_4x8_q8_0 #elif defined(__s390x__) +// PTQ1_0 currently has only the generic vec_dot; alias it here until a SIMD version lands +#define ggml_vec_dot_ptq1_0_q8_0_generic ggml_vec_dot_ptq1_0_q8_0 // quants.c #define quantize_row_q8_K_generic quantize_row_q8_K #define ggml_vec_dot_nvfp4_q8_0_generic ggml_vec_dot_nvfp4_q8_0 @@ -336,6 +349,8 @@ #define ggml_gemm_q1_0_4x8_q8_0_generic ggml_gemm_q1_0_4x8_q8_0 #define ggml_gemm_pq2_0_4x8_q8_0_generic ggml_gemm_pq2_0_4x8_q8_0 #elif defined(__wasm__) +// PTQ1_0 currently has only the generic vec_dot; alias it here until a SIMD version lands +#define ggml_vec_dot_ptq1_0_q8_0_generic ggml_vec_dot_ptq1_0_q8_0 // quants.c #define ggml_vec_dot_tq1_0_q8_K_generic ggml_vec_dot_tq1_0_q8_K #define ggml_vec_dot_tq2_0_q8_K_generic ggml_vec_dot_tq2_0_q8_K diff --git a/ggml/src/ggml-cpu/ggml-cpu.c b/ggml/src/ggml-cpu/ggml-cpu.c index 79c49c5e4d3b..7a1bcdf355b3 100644 --- a/ggml/src/ggml-cpu/ggml-cpu.c +++ b/ggml/src/ggml-cpu/ggml-cpu.c @@ -242,6 +242,12 @@ static const struct ggml_type_traits_cpu type_traits_cpu[GGML_TYPE_COUNT] = { .vec_dot_type = GGML_TYPE_Q8_0, .nrows = 1, }, + [GGML_TYPE_PTQ1_0] = { + .from_float = quantize_row_ptq1_0, + .vec_dot = ggml_vec_dot_ptq1_0_q8_0, + .vec_dot_type = GGML_TYPE_Q8_0, + .nrows = 1, + }, [GGML_TYPE_Q4_0] = { .from_float = quantize_row_q4_0, .vec_dot = ggml_vec_dot_q4_0_q8_0, diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp index f5bcbfb6389e..35697d0787d7 100644 --- a/ggml/src/ggml-cpu/ops.cpp +++ b/ggml/src/ggml-cpu/ops.cpp @@ -667,6 +667,7 @@ void ggml_compute_forward_add( case GGML_TYPE_Q1_0: case GGML_TYPE_Q2_0: case GGML_TYPE_PQ2_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: @@ -1119,6 +1120,7 @@ void ggml_compute_forward_add1( case GGML_TYPE_Q1_0: case GGML_TYPE_Q2_0: case GGML_TYPE_PQ2_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: @@ -1251,6 +1253,7 @@ void ggml_compute_forward_acc( case GGML_TYPE_Q1_0: case GGML_TYPE_Q2_0: case GGML_TYPE_PQ2_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: @@ -4522,6 +4525,7 @@ void ggml_compute_forward_out_prod( case GGML_TYPE_Q1_0: case GGML_TYPE_Q2_0: case GGML_TYPE_PQ2_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: @@ -4799,6 +4803,7 @@ void ggml_compute_forward_set( case GGML_TYPE_Q1_0: case GGML_TYPE_Q2_0: case GGML_TYPE_PQ2_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: @@ -5025,6 +5030,7 @@ void ggml_compute_forward_get_rows( case GGML_TYPE_Q1_0: case GGML_TYPE_Q2_0: case GGML_TYPE_PQ2_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: @@ -5783,6 +5789,7 @@ void ggml_compute_forward_clamp( case GGML_TYPE_Q1_0: case GGML_TYPE_Q2_0: case GGML_TYPE_PQ2_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: diff --git a/ggml/src/ggml-cpu/quants.c b/ggml/src/ggml-cpu/quants.c index 1971d57a0224..3e275dc70bf3 100644 --- a/ggml/src/ggml-cpu/quants.c +++ b/ggml/src/ggml-cpu/quants.c @@ -34,6 +34,10 @@ void quantize_row_pq2_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, i quantize_row_pq2_0_ref(x, y, k); } +void quantize_row_ptq1_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { + quantize_row_ptq1_0_ref(x, y, k); +} + void quantize_row_q4_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { quantize_row_q4_0_ref(x, y, k); } @@ -274,6 +278,74 @@ void ggml_vec_dot_pq2_0_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, *s = sumf; } +// PTQ1_0 x Q8_0. The trits are stored base-3 interleaved rather than in element +// order, so decode a block into element order first using the same traversal as +// dequantize_row_ptq1_0 -- that keeps the two provably in step. Four Q8_0 blocks +// cover one 128-wide PTQ1_0 block. +void ggml_vec_dot_ptq1_0_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc) { + const int qk = QK_PTQ1_0; + const int nb = n / qk; + + assert(n % qk == 0); + assert(nrc == 1); + UNUSED(nrc); + UNUSED(bx); + UNUSED(by); + UNUSED(bs); + + const block_ptq1_0 * GGML_RESTRICT x = vx; + const block_q8_0 * GGML_RESTRICT y = vy; + + static const uint8_t pow3[6] = {1, 3, 9, 27, 81, 243}; + static const size_t stages[3] = {32, 16, 8}; + + float sumf = 0.0f; + + for (int i = 0; i < nb; i++) { + int8_t q[QK_PTQ1_0]; + int o = 0; + + size_t j = 0; + for (size_t st = 0; st < 3; ++st) { + const size_t c = stages[st]; + for (; j + c <= sizeof(x->qs); j += c) { + for (size_t nn = 0; nn < 5; ++nn) { + for (size_t m = 0; m < c; ++m) { + const uint8_t v = x[i].qs[j + m] * pow3[nn]; + const int16_t xi = ((uint16_t) v * 3) >> 8; + q[o++] = (int8_t) (xi - 1); + } + } + } + } + for (size_t nn = 0; nn < 4; ++nn) { + for (size_t h = 0; h < sizeof(x->qh); ++h) { + const uint8_t v = x[i].qh[h] * pow3[nn]; + const int16_t xi = ((uint16_t) v * 3) >> 8; + q[o++] = (int8_t) (xi - 1); + } + } + assert(o == QK_PTQ1_0); + + const float d0 = GGML_CPU_FP16_TO_FP32(x[i].d); + float sumi = 0.0f; + + for (int k = 0; k < 4; k++) { + const block_q8_0 * GGML_RESTRICT yb = &y[i * 4 + k]; + const float d1 = GGML_CPU_FP16_TO_FP32(yb->d); + int sumi_block = 0; + for (int b = 0; b < 32; ++b) { + sumi_block += (int) q[k*32 + b] * (int) yb->qs[b]; + } + sumi += d1 * sumi_block; + } + + sumf += d0 * sumi; + } + + *s = sumf; +} + void ggml_vec_dot_q4_0_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc) { const int qk = QK8_0; const int nb = n / qk; diff --git a/ggml/src/ggml-cpu/quants.h b/ggml/src/ggml-cpu/quants.h index 2cac90c6b0a0..f00f35c3fe85 100644 --- a/ggml/src/ggml-cpu/quants.h +++ b/ggml/src/ggml-cpu/quants.h @@ -15,6 +15,7 @@ extern "C" { void quantize_row_q1_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); void quantize_row_q2_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); void quantize_row_pq2_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +void quantize_row_ptq1_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); void quantize_row_q4_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); void quantize_row_q4_1(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); void quantize_row_q5_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); @@ -42,6 +43,7 @@ void quantize_row_iq4_xs (const float * GGML_RESTRICT x, void * GGML_RESTRICT y, void ggml_vec_dot_q1_0_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); void ggml_vec_dot_q2_0_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); void ggml_vec_dot_pq2_0_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); +void ggml_vec_dot_ptq1_0_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); void ggml_vec_dot_q4_0_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); void ggml_vec_dot_q4_1_q8_1(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); void ggml_vec_dot_q5_0_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); @@ -76,6 +78,7 @@ void quantize_row_q8_1_generic(const float * GGML_RESTRICT x, void * GGML_RESTRI void quantize_row_q8_K_generic(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); void ggml_vec_dot_q1_0_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); void ggml_vec_dot_pq2_0_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); +void ggml_vec_dot_ptq1_0_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); void ggml_vec_dot_q2_0_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); void ggml_vec_dot_q4_0_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); void ggml_vec_dot_q4_1_q8_1_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index b2d82beaf7c8..ef929d3d7842 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -992,6 +992,40 @@ struct ggml_cuda_type_traits { static constexpr int qi = QI_PQ2_0; }; +// PTQ1_0 packs trits base-3, five per byte, so element order is not positional. It +// follows the CPU codec exactly: a 16-byte chunk of qs, then an 8-byte chunk, then qh +// at four trits per byte. Trits come out by the base-3 remainder recurrence +// t = (v*3)>>8 with v = (v*3)&0xFF, two integer ops per step and no table. +static __device__ __forceinline__ int ptq1_0_trit(const block_ptq1_0 * x, const int e) { + uint8_t b; + int n; + if (e < 80) { // qs[0..15], chunk of 16 + b = x->qs[e & 15]; n = e >> 4; + } else if (e < 120) { // qs[16..23], chunk of 8 + const int t = e - 80; + b = x->qs[16 + (t & 7)]; n = t >> 3; + } else { // qh[0..1], four trits per byte + const int t = e - 120; + b = x->qh[t & 1]; n = t >> 1; + } + + uint32_t v = b; +#pragma unroll + for (int i = 0; i < 4; ++i) { + if (i < n) { + v = (v * 3) & 0xFF; + } + } + return (int) ((v * 3) >> 8) - 1; +} + +template<> +struct ggml_cuda_type_traits { + static constexpr int qk = QK_PTQ1_0; + static constexpr int qr = QR_PTQ1_0; + static constexpr int qi = QI_PTQ1_0; +}; + template<> struct ggml_cuda_type_traits { static constexpr int qk = QK4_0; diff --git a/ggml/src/ggml-cuda/convert.cu b/ggml/src/ggml-cuda/convert.cu index 807f2caa73ff..f271c00912e9 100644 --- a/ggml/src/ggml-cuda/convert.cu +++ b/ggml/src/ggml-cuda/convert.cu @@ -463,6 +463,8 @@ to_bf16_cuda_t ggml_get_to_bf16_cuda(ggml_type type) { return dequantize_block_cont_cuda; case GGML_TYPE_PQ2_0: return dequantize_block_cont_cuda; + case GGML_TYPE_PTQ1_0: + return dequantize_block_cont_cuda; case GGML_TYPE_Q4_0: return dequantize_row_q4_0_cuda; case GGML_TYPE_Q4_1: @@ -522,6 +524,8 @@ to_fp16_cuda_t ggml_get_to_fp16_cuda(ggml_type type) { return dequantize_block_cont_cuda; case GGML_TYPE_PQ2_0: return dequantize_block_cont_cuda; + case GGML_TYPE_PTQ1_0: + return dequantize_block_cont_cuda; case GGML_TYPE_Q4_0: return dequantize_row_q4_0_cuda; case GGML_TYPE_Q4_1: @@ -584,6 +588,8 @@ to_fp32_cuda_t ggml_get_to_fp32_cuda(ggml_type type) { return dequantize_block_cont_cuda; case GGML_TYPE_PQ2_0: return dequantize_block_cont_cuda; + case GGML_TYPE_PTQ1_0: + return dequantize_block_cont_cuda; case GGML_TYPE_Q4_0: return dequantize_row_q4_0_cuda; case GGML_TYPE_Q4_1: @@ -645,6 +651,8 @@ to_fp16_nc_cuda_t ggml_get_to_fp16_nc_cuda(ggml_type type) { return dequantize_block_cuda; case GGML_TYPE_PQ2_0: return dequantize_block_cuda; + case GGML_TYPE_PTQ1_0: + return dequantize_block_cuda; case GGML_TYPE_Q4_0: return dequantize_block_cuda; case GGML_TYPE_Q4_1: @@ -672,6 +680,8 @@ to_bf16_nc_cuda_t ggml_get_to_bf16_nc_cuda(ggml_type type) { return dequantize_block_cuda; case GGML_TYPE_PQ2_0: return dequantize_block_cuda; + case GGML_TYPE_PTQ1_0: + return dequantize_block_cuda; case GGML_TYPE_Q4_0: return dequantize_block_cuda; case GGML_TYPE_Q4_1: @@ -699,6 +709,8 @@ to_fp32_nc_cuda_t ggml_get_to_fp32_nc_cuda(ggml_type type) { return dequantize_block_cuda; case GGML_TYPE_PQ2_0: return dequantize_block_cuda; + case GGML_TYPE_PTQ1_0: + return dequantize_block_cuda; case GGML_TYPE_Q4_0: return dequantize_block_cuda; case GGML_TYPE_Q4_1: diff --git a/ggml/src/ggml-cuda/dequantize.cuh b/ggml/src/ggml-cuda/dequantize.cuh index e3f130f1c8ab..7a1deec8b829 100644 --- a/ggml/src/ggml-cuda/dequantize.cuh +++ b/ggml/src/ggml-cuda/dequantize.cuh @@ -43,6 +43,16 @@ static __device__ __forceinline__ void dequantize_q2_0(const void * vx, const in v.y = (c1 - 1) * d; } + +static __device__ __forceinline__ void dequantize_ptq1_0(const void * vx, const int64_t ib, const int iqs, float2 & v){ + const block_ptq1_0 * x = (const block_ptq1_0 *) vx; + + const float d = x[ib].d; + + v.x = ptq1_0_trit(&x[ib], iqs + 0) * d; + v.y = ptq1_0_trit(&x[ib], iqs + 1) * d; +} + static __device__ __forceinline__ void dequantize_pq2_0(const void * vx, const int64_t ib, const int iqs, float2 & v){ const block_pq2_0 * x = (const block_pq2_0 *) vx; diff --git a/ggml/src/ggml-cuda/getrows.cu b/ggml/src/ggml-cuda/getrows.cu index 49f5bd28dfac..6b4453bc69ce 100644 --- a/ggml/src/ggml-cuda/getrows.cu +++ b/ggml/src/ggml-cuda/getrows.cu @@ -324,6 +324,10 @@ static void ggml_cuda_get_rows_switch_src0_type( get_rows_cuda_q(src0_d, src1_d, dst_d, ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream); break; + case GGML_TYPE_PTQ1_0: + get_rows_cuda_q(src0_d, src1_d, dst_d, + ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream); + break; case GGML_TYPE_Q2_0: get_rows_cuda_q(src0_d, src1_d, dst_d, ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream); diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 9178d32fe717..dedf60b30d80 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -4988,6 +4988,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g case GGML_TYPE_Q1_0: case GGML_TYPE_Q2_0: case GGML_TYPE_PQ2_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: @@ -5028,6 +5029,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g case GGML_TYPE_Q1_0: case GGML_TYPE_Q2_0: case GGML_TYPE_PQ2_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index 768b2394abc3..ede8a65e0154 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -13,6 +13,7 @@ static constexpr __device__ vec_dot_q_cuda_t get_vec_dot_q_cuda(ggml_type type) case GGML_TYPE_Q1_0: return vec_dot_q1_0_q8_1; case GGML_TYPE_Q2_0: return vec_dot_q2_0_q8_1; case GGML_TYPE_PQ2_0: return vec_dot_pq2_0_q8_1; + case GGML_TYPE_PTQ1_0: return vec_dot_ptq1_0_q8_1; case GGML_TYPE_Q4_0: return vec_dot_q4_0_q8_1; case GGML_TYPE_Q4_1: return vec_dot_q4_1_q8_1; case GGML_TYPE_Q5_0: return vec_dot_q5_0_q8_1; @@ -43,6 +44,7 @@ static constexpr __host__ __device__ int get_vdr_mmvq(ggml_type type) { case GGML_TYPE_Q1_0: return VDR_Q1_0_Q8_1_MMVQ; case GGML_TYPE_Q2_0: return VDR_Q2_0_Q8_1_MMVQ; case GGML_TYPE_PQ2_0: return VDR_PQ2_0_Q8_1_MMVQ; + case GGML_TYPE_PTQ1_0: return VDR_PTQ1_0_Q8_1_MMVQ; case GGML_TYPE_Q4_0: return VDR_Q4_0_Q8_1_MMVQ; case GGML_TYPE_Q4_1: return VDR_Q4_1_Q8_1_MMVQ; case GGML_TYPE_Q5_0: return VDR_Q5_0_Q8_1_MMVQ; @@ -663,24 +665,52 @@ static __global__ void mul_mat_vec_q( // x block quant index when casting the quants to int const int kqs = vdr * (tid % (qi/vdr)); -#pragma unroll - for (int j = 0; j < ncols_dst; ++j) { -#pragma unroll +#if !defined(GGML_USE_HIP) + if constexpr (type == GGML_TYPE_PTQ1_0 && ncols_dst > 1 && ncols_dst <= 3) { +# pragma unroll for (int i = 0; i < rows_per_cuda_block; ++i) { - tmp[j][i] += vec_dot_q_cuda( - vx, &y[j*stride_col_y + kby], kbx_offset + i*stride_row_x + kbx, kqs); + float dots[ncols_dst]; + vec_dot_ptq1_0_q8_1_multi(vx, &y[kby], kbx_offset + i * stride_row_x + kbx, kqs, + stride_col_y, dots); +# pragma unroll + for (int j = 0; j < ncols_dst; ++j) { + tmp[j][i] += dots[j]; + } + if constexpr (has_fusion) { if (use_gate) { - tmp_gate[j][i] += vec_dot_q_cuda( - vgate, &y[j*stride_col_y + kby], kbx_offset + i*stride_row_x + kbx, kqs); + vec_dot_ptq1_0_q8_1_multi(vgate, &y[kby], kbx_offset + i * stride_row_x + kbx, kqs, + stride_col_y, dots); +# pragma unroll + for (int j = 0; j < ncols_dst; ++j) { + tmp_gate[j][i] += dots[j]; + } + } + } + } + } else +#endif + { +#pragma unroll + for (int j = 0; j < ncols_dst; ++j) { +#pragma unroll + for (int i = 0; i < rows_per_cuda_block; ++i) { + tmp[j][i] += + vec_dot_q_cuda(vx, &y[j * stride_col_y + kby], kbx_offset + i * stride_row_x + kbx, kqs); + if constexpr (has_fusion) { + if (use_gate) { + tmp_gate[j][i] += vec_dot_q_cuda(vgate, &y[j * stride_col_y + kby], + kbx_offset + i * stride_row_x + kbx, kqs); + } } } } } } - __shared__ float tmp_shared[nwarps-1 > 0 ? nwarps-1 : 1][ncols_dst][rows_per_cuda_block][warp_size]; - [[maybe_unused]] __shared__ float tmp_shared_gate[(has_fusion && (nwarps-1 > 0)) ? nwarps-1 : 1][ncols_dst][rows_per_cuda_block][warp_size]; + __shared__ float tmp_shared[nwarps - 1 > 0 ? nwarps - 1 : 1][ncols_dst][rows_per_cuda_block][warp_size]; + [[maybe_unused]] __shared__ float tmp_shared_gate[(has_fusion && (nwarps - 1 > 0)) ? nwarps - 1 : 1][ncols_dst] + [rows_per_cuda_block][warp_size]; if (threadIdx.y > 0) { #pragma unroll @@ -1120,6 +1150,12 @@ static void mul_mat_vec_q_switch_type( nchannels_x, nchannels_y, nchannels_dst, stride_channel_x, stride_channel_y, stride_channel_dst, nsamples_x, nsamples_dst, stride_sample_x, stride_sample_y, stride_sample_dst, ids_stride, stream); break; + case GGML_TYPE_PTQ1_0: + mul_mat_vec_q_switch_ncols_dst + (vx, vy, ids, fusion, dst, ncols_x, nrows_x, ncols_dst, stride_row_x, stride_col_y, stride_col_dst, + nchannels_x, nchannels_y, nchannels_dst, stride_channel_x, stride_channel_y, stride_channel_dst, + nsamples_x, nsamples_dst, stride_sample_x, stride_sample_y, stride_sample_dst, ids_stride, stream); + break; case GGML_TYPE_Q2_0: mul_mat_vec_q_switch_ncols_dst (vx, vy, ids, fusion, dst, ncols_x, nrows_x, ncols_dst, stride_row_x, stride_col_y, stride_col_dst, diff --git a/ggml/src/ggml-cuda/vecdotq.cuh b/ggml/src/ggml-cuda/vecdotq.cuh index f6b1acd591c2..fa6cdabd5d84 100644 --- a/ggml/src/ggml-cuda/vecdotq.cuh +++ b/ggml/src/ggml-cuda/vecdotq.cuh @@ -113,6 +113,7 @@ static __device__ __forceinline__ uint32_t unpack_ksigns(const uint8_t v) { #define VDR_Q2_0_Q8_1_MMQ 2 // Q2_0 group 64: 128 bits (4 ints) per block, 2 32-element chunks #define VDR_PQ2_0_Q8_1_MMVQ 1 // one 32-element chunk at a time (same per-chunk codec as Q2_0) +#define VDR_PTQ1_0_Q8_1_MMVQ 4 // whole 128 block per call: keeps the byte walk uniform across lanes #define VDR_PQ2_0_Q8_1_MMQ 2 // Q2_0 group 128: 4 32-element chunks per block #define VDR_Q4_0_Q8_1_MMVQ 2 @@ -802,9 +803,158 @@ static __device__ __forceinline__ float vec_dot_q2_0_q8_1( return d2 * d8 * sumi; } -static __device__ __forceinline__ float vec_dot_pq2_0_q8_1( - const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) { +#if !defined(GGML_USE_HIP) +template +static __device__ __forceinline__ void vec_dot_ptq1_0_q8_1_multi(const void * __restrict__ vbq, + const block_q8_1 * __restrict__ bq8_1, + const int & kbx, + const int & iqs, + const uint32_t stride_col_y, + float * result) { + const block_ptq1_0 * bq = (const block_ptq1_0 *) vbq + kbx; + int sumi[ncols_dst][4] = {}; + + // Widen four bytes to 16-bit lanes so multiply-by-three cannot carry between bytes. +# pragma unroll + for (int g = 0; g < 4; ++g) { + const uint32_t packed = get_int_b4(bq->qs, g); + uint32_t v_lo = __byte_perm(packed, 0, 0x4140); + uint32_t v_hi = __byte_perm(packed, 0, 0x4342); + +# pragma unroll + for (int t = 0; t < 5; ++t) { + const uint32_t w_lo = v_lo * 3; + const uint32_t w_hi = v_hi * 3; + v_lo = w_lo & 0x00FF00FF; + v_hi = w_hi & 0x00FF00FF; + + const int q = __vsub4(__byte_perm(w_lo, w_hi, 0x7531), 0x01010101); + const int e = t * 16 + 4 * g; +# pragma unroll + for (int j = 0; j < ncols_dst; ++j) { + const int u = get_int_b4(bq8_1[j * stride_col_y + iqs + (e >> 5)].qs, (e & 31) >> 2); + sumi[j][e >> 5] = ggml_cuda_dp4a(q, u, sumi[j][e >> 5]); + } + } + } + +# pragma unroll + for (int g = 0; g < 2; ++g) { + const uint32_t packed = get_int_b4(bq->qs + 16, g); + uint32_t v_lo = __byte_perm(packed, 0, 0x4140); + uint32_t v_hi = __byte_perm(packed, 0, 0x4342); + +# pragma unroll + for (int t = 0; t < 5; ++t) { + const uint32_t w_lo = v_lo * 3; + const uint32_t w_hi = v_hi * 3; + v_lo = w_lo & 0x00FF00FF; + v_hi = w_hi & 0x00FF00FF; + + const int q = __vsub4(__byte_perm(w_lo, w_hi, 0x7531), 0x01010101); + const int e = 80 + t * 8 + 4 * g; +# pragma unroll + for (int j = 0; j < ncols_dst; ++j) { + const int u = get_int_b4(bq8_1[j * stride_col_y + iqs + (e >> 5)].qs, (e & 31) >> 2); + sumi[j][e >> 5] = ggml_cuda_dp4a(q, u, sumi[j][e >> 5]); + } + } + } + + uint32_t v = (uint32_t) bq->qh[0] | ((uint32_t) bq->qh[1] << 16); +# pragma unroll + for (int t = 0; t < 4; t += 2) { + const uint32_t w0 = v * 3; + v = w0 & 0x00FF00FF; + const uint32_t w1 = v * 3; + v = w1 & 0x00FF00FF; + + const int q = __vsub4(__byte_perm(w0, w1, 0x7531), 0x01010101); +# pragma unroll + for (int j = 0; j < ncols_dst; ++j) { + const int u = get_int_b4(bq8_1[j * stride_col_y + iqs + 3].qs, 6 + t / 2); + sumi[j][3] = ggml_cuda_dp4a(q, u, sumi[j][3]); + } + } + + const float d = (float) bq->d; +# pragma unroll + for (int j = 0; j < ncols_dst; ++j) { + float acc = 0.0f; +# pragma unroll + for (int k = 0; k < 4; ++k) { + acc += __low2float(bq8_1[j * stride_col_y + iqs + k].ds) * (float) sumi[j][k]; + } + result[j] = d * acc; + } +} +#endif + +// PTQ1_0 x Q8_1. One call consumes the full 128-weight block. +static __device__ __forceinline__ float vec_dot_ptq1_0_q8_1(const void * __restrict__ vbq, + const block_q8_1 * __restrict__ bq8_1, + const int & kbx, + const int & iqs) { +#if defined(GGML_USE_HIP) + const block_ptq1_0 * bq = (const block_ptq1_0 *) vbq + kbx; + int sumi[4] = { 0, 0, 0, 0 }; + +# pragma unroll + for (int m = 0; m < 16; ++m) { + uint32_t v = bq->qs[m]; +# pragma unroll + for (int t = 0; t < 5; ++t) { + const uint32_t w = v * 3; + const int q = (int) (w >> 8) - 1; + v = w & 0xFF; + const int e = t * 16 + m; + sumi[e >> 5] += q * (int) bq8_1[iqs + (e >> 5)].qs[e & 31]; + } + } + +# pragma unroll + for (int m = 0; m < 8; ++m) { + uint32_t v = bq->qs[16 + m]; +# pragma unroll + for (int t = 0; t < 5; ++t) { + const uint32_t w = v * 3; + const int q = (int) (w >> 8) - 1; + v = w & 0xFF; + const int e = 80 + t * 8 + m; + sumi[e >> 5] += q * (int) bq8_1[iqs + (e >> 5)].qs[e & 31]; + } + } + +# pragma unroll + for (int h = 0; h < 2; ++h) { + uint32_t v = bq->qh[h]; +# pragma unroll + for (int t = 0; t < 4; ++t) { + const uint32_t w = v * 3; + const int q = (int) (w >> 8) - 1; + v = w & 0xFF; + const int e = 120 + t * 2 + h; + sumi[e >> 5] += q * (int) bq8_1[iqs + (e >> 5)].qs[e & 31]; + } + } + + float acc = 0.0f; +# pragma unroll + for (int k = 0; k < 4; ++k) { + acc += __low2float(bq8_1[iqs + k].ds) * (float) sumi[k]; + } + return (float) bq->d * acc; +#else + float result; + vec_dot_ptq1_0_q8_1_multi<1>(vbq, bq8_1, kbx, iqs, 0, &result); + return result; +#endif +} +static __device__ __forceinline__ float vec_dot_pq2_0_q8_1(const void * __restrict__ vbq, + const block_q8_1 * __restrict__ bq8_1, + const int & kbx, + const int & iqs) { const block_pq2_0 * bq2_0 = (const block_pq2_0 *) vbq + kbx; // Q2_0 group 128: 128 elements, ONE scale, processed as four 32-element chunks diff --git a/ggml/src/ggml-metal/ggml-metal-device.cpp b/ggml/src/ggml-metal/ggml-metal-device.cpp index 4870bdcb968d..e85c76e0cca7 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.cpp +++ b/ggml/src/ggml-metal/ggml-metal-device.cpp @@ -871,6 +871,11 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv(ggml_meta nsg = N_SG_PQ2_0; nr0 = N_R0_PQ2_0; } break; + case GGML_TYPE_PTQ1_0: + { + nsg = N_SG_PTQ1_0; + nr0 = N_R0_PTQ1_0; + } break; case GGML_TYPE_Q4_0: { nsg = N_SG_Q4_0; @@ -1110,6 +1115,11 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_id(ggml_m nsg = N_SG_PQ2_0; nr0 = N_R0_PQ2_0; } break; + case GGML_TYPE_PTQ1_0: + { + nsg = N_SG_PTQ1_0; + nr0 = N_R0_PTQ1_0; + } break; case GGML_TYPE_Q4_0: { nsg = N_SG_Q4_0; diff --git a/ggml/src/ggml-metal/ggml-metal-device.m b/ggml/src/ggml-metal/ggml-metal-device.m index 08c0ef6ce67d..43b1d8c6f306 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.m +++ b/ggml/src/ggml-metal/ggml-metal-device.m @@ -1732,6 +1732,7 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te case GGML_TYPE_Q1_0: case GGML_TYPE_Q2_0: case GGML_TYPE_PQ2_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: @@ -1762,6 +1763,7 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te case GGML_TYPE_Q1_0: case GGML_TYPE_Q2_0: case GGML_TYPE_PQ2_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: diff --git a/ggml/src/ggml-metal/ggml-metal-impl.h b/ggml/src/ggml-metal/ggml-metal-impl.h index 836cbbcb5c9e..1151e5099fc8 100644 --- a/ggml/src/ggml-metal/ggml-metal-impl.h +++ b/ggml/src/ggml-metal/ggml-metal-impl.h @@ -29,6 +29,8 @@ #define N_R0_PQ2_0 8 #define N_SG_PQ2_0 2 +#define N_R0_PTQ1_0 4 +#define N_SG_PTQ1_0 1 #define N_R0_Q4_0 4 #define N_SG_Q4_0 2 diff --git a/ggml/src/ggml-metal/ggml-metal-ops.cpp b/ggml/src/ggml-metal/ggml-metal-ops.cpp index 6be0f52015b2..38744c2d3cdd 100644 --- a/ggml/src/ggml-metal/ggml-metal-ops.cpp +++ b/ggml/src/ggml-metal/ggml-metal-ops.cpp @@ -2565,6 +2565,7 @@ int ggml_metal_op_mul_mat(ggml_metal_op_t ctx, int idx) { op->src[0]->type == GGML_TYPE_Q1_0 || op->src[0]->type == GGML_TYPE_Q2_0 || op->src[0]->type == GGML_TYPE_PQ2_0 || + op->src[0]->type == GGML_TYPE_PTQ1_0 || op->src[0]->type == GGML_TYPE_Q4_0 || op->src[0]->type == GGML_TYPE_Q4_1 || op->src[0]->type == GGML_TYPE_Q5_0 || diff --git a/ggml/src/ggml-metal/kernels/dequantize.h b/ggml/src/ggml-metal/kernels/dequantize.h index 1254d190c9b1..1e01dcd81d9e 100644 --- a/ggml/src/ggml-metal/kernels/dequantize.h +++ b/ggml/src/ggml-metal/kernels/dequantize.h @@ -161,6 +161,89 @@ void dequantize_pq2_0_t4(device const block_pq2_0 * xb, short il, thread type4 & reg = (type4) reg_f; } +// PTQ1_0: one lookup replaces the base-3 arithmetic. Entry b holds the five trits of +// byte b packed two bits each, low trit first, as 0/1/2 (subtract 1 for the value). +// Generated from, and verified against, the reference decode ((b*3^n & 0xFF)*3)>>8 for +// all 256 bytes and all five trit positions. This puts extraction on par with a 2-bit +// format: shift, mask, subtract, with a single cached table read amortised over 5 trits. +constant ushort ptq1_0_lut[256] = { + 0, 0, 256, 512, 64, 320, 576, 128, 384, 640, 16, 272, 528, 80, 336, 592, + 144, 400, 656, 32, 32, 288, 544, 96, 352, 608, 160, 416, 672, 4, 260, 516, + 68, 324, 580, 132, 388, 644, 20, 276, 276, 532, 84, 340, 596, 148, 404, 660, + 36, 292, 548, 100, 356, 612, 164, 420, 676, 8, 264, 520, 520, 72, 328, 584, + 136, 392, 648, 24, 280, 536, 88, 344, 600, 152, 408, 664, 40, 296, 552, 552, + 104, 360, 616, 168, 424, 680, 1, 257, 513, 65, 321, 577, 129, 385, 641, 17, + 273, 529, 81, 81, 337, 593, 145, 401, 657, 33, 289, 545, 97, 353, 609, 161, + 417, 673, 5, 261, 517, 69, 325, 325, 581, 133, 389, 645, 21, 277, 533, 85, + 341, 597, 149, 405, 661, 37, 293, 549, 101, 357, 357, 613, 165, 421, 677, 9, + 265, 521, 73, 329, 585, 137, 393, 649, 25, 281, 537, 89, 345, 601, 601, 153, + 409, 665, 41, 297, 553, 105, 361, 617, 169, 425, 681, 2, 258, 514, 66, 322, + 578, 130, 130, 386, 642, 18, 274, 530, 82, 338, 594, 146, 402, 658, 34, 290, + 546, 98, 354, 610, 162, 162, 418, 674, 6, 262, 518, 70, 326, 582, 134, 390, + 646, 22, 278, 534, 86, 342, 598, 150, 406, 406, 662, 38, 294, 550, 102, 358, + 614, 166, 422, 678, 10, 266, 522, 74, 330, 586, 138, 394, 650, 650, 26, 282, + 538, 90, 346, 602, 154, 410, 666, 42, 298, 554, 106, 362, 618, 170, 426, 682, +}; + +// PTQ1_0 stores trits base-3 packed, so element order is not positional: it follows +// the CPU codec's 16-then-8 byte staging over qs, then qh at 4 trits per byte. +// Map an element index to its byte and trit rather than assuming contiguity. +inline float ptq1_0_elem(device const block_ptq1_0 * xb, int e) { + uchar b; + short n; + if (e < 80) { // qs[0..15], chunk of 16, 5 trits per byte + b = xb->qs[e & 15]; n = e >> 4; + } else if (e < 120) { // qs[16..23], chunk of 8 + const int t = e - 80; + b = xb->qs[16 + (t & 7)]; n = t >> 3; + } else { // qh[0..1], 4 trits per byte + const int t = e - 120; + b = xb->qh[t & 1]; n = t >> 1; + } + return (float) ((int) ((ptq1_0_lut[b] >> (2*n)) & 3) - 1); +} + +template +void dequantize_ptq1_0(device const block_ptq1_0 * xb, short il, thread type4x4 & reg) { + const float d = xb->d; + + float4x4 reg_f; + + // il 0..4 addresses the 16-byte chunk at a single trit index, so the pow3 factor + // and the branch hoist out and this becomes 16 plain byte loads. il 5..7 straddle + // trit indices, or reach into qh, and take the general accessor. + if (il < 5) { + const short sh = 2*il; + device const uchar * qs = xb->qs; + for (int i = 0; i < 4; i++) { + for (int j = 0; j < 4; j++) { + reg_f[i][j] = (float) ((int) ((ptq1_0_lut[qs[i*4 + j]] >> sh) & 3) - 1) * d; + } + } + } else { + const int base = il * 16; + for (int i = 0; i < 4; i++) { + for (int j = 0; j < 4; j++) { + reg_f[i][j] = ptq1_0_elem(xb, base + i*4 + j) * d; + } + } + } + + reg = (type4x4) reg_f; +} + +template +void dequantize_ptq1_0_t4(device const block_ptq1_0 * xb, short il, thread type4 & reg) { + const float d = xb->d; + const int base = il * 4; + + float4 reg_f; + for (int j = 0; j < 4; j++) { + reg_f[j] = ptq1_0_elem(xb, base + j) * d; + } + reg = (type4) reg_f; +} + template void dequantize_q4_0(device const block_q4_0 * xb, short il, thread type4x4 & reg) { device const uint16_t * qs = ((device const uint16_t *)xb + 1); diff --git a/ggml/src/ggml-metal/kernels/mul_mm.metal b/ggml/src/ggml-metal/kernels/mul_mm.metal index 443f41d52332..b0029fdc6ad2 100644 --- a/ggml/src/ggml-metal/kernels/mul_mm.metal +++ b/ggml/src/ggml-metal/kernels/mul_mm.metal @@ -744,6 +744,7 @@ template [[host_name("kernel_mul_mm_bf16_f32")]] kernel mul_mm_t kernel_mul_m template [[host_name("kernel_mul_mm_q1_0_f32")]] kernel mul_mm_t kernel_mul_mm; template [[host_name("kernel_mul_mm_q2_0_f32")]] kernel mul_mm_t kernel_mul_mm; template [[host_name("kernel_mul_mm_pq2_0_f32")]] kernel mul_mm_t kernel_mul_mm; +template [[host_name("kernel_mul_mm_ptq1_0_f32")]] kernel mul_mm_t kernel_mul_mm; template [[host_name("kernel_mul_mm_q4_0_f32")]] kernel mul_mm_t kernel_mul_mm; template [[host_name("kernel_mul_mm_q4_1_f32")]] kernel mul_mm_t kernel_mul_mm; template [[host_name("kernel_mul_mm_q5_0_f32")]] kernel mul_mm_t kernel_mul_mm; @@ -771,6 +772,7 @@ template [[host_name("kernel_mul_mm_f16_f16")]] kernel mul_mm_t kernel_mul_m template [[host_name("kernel_mul_mm_q1_0_f16")]] kernel mul_mm_t kernel_mul_mm; template [[host_name("kernel_mul_mm_q2_0_f16")]] kernel mul_mm_t kernel_mul_mm; template [[host_name("kernel_mul_mm_pq2_0_f16")]] kernel mul_mm_t kernel_mul_mm; +template [[host_name("kernel_mul_mm_ptq1_0_f16")]] kernel mul_mm_t kernel_mul_mm; template [[host_name("kernel_mul_mm_q4_0_f16")]] kernel mul_mm_t kernel_mul_mm; template [[host_name("kernel_mul_mm_q4_1_f16")]] kernel mul_mm_t kernel_mul_mm; template [[host_name("kernel_mul_mm_q5_0_f16")]] kernel mul_mm_t kernel_mul_mm; @@ -807,6 +809,7 @@ template [[host_name("kernel_mul_mm_id_bf16_f32")]] kernel mul_mm_id kernel_m template [[host_name("kernel_mul_mm_id_q1_0_f32")]] kernel mul_mm_id kernel_mul_mm_id; template [[host_name("kernel_mul_mm_id_q2_0_f32")]] kernel mul_mm_id kernel_mul_mm_id; template [[host_name("kernel_mul_mm_id_pq2_0_f32")]] kernel mul_mm_id kernel_mul_mm_id; +template [[host_name("kernel_mul_mm_id_ptq1_0_f32")]] kernel mul_mm_id kernel_mul_mm_id; template [[host_name("kernel_mul_mm_id_q4_0_f32")]] kernel mul_mm_id kernel_mul_mm_id; template [[host_name("kernel_mul_mm_id_q4_1_f32")]] kernel mul_mm_id kernel_mul_mm_id; template [[host_name("kernel_mul_mm_id_q5_0_f32")]] kernel mul_mm_id kernel_mul_mm_id; @@ -834,6 +837,7 @@ template [[host_name("kernel_mul_mm_id_f16_f16")]] kernel mul_mm_id kernel_m template [[host_name("kernel_mul_mm_id_q1_0_f16")]] kernel mul_mm_id kernel_mul_mm_id; template [[host_name("kernel_mul_mm_id_q2_0_f16")]] kernel mul_mm_id kernel_mul_mm_id; template [[host_name("kernel_mul_mm_id_pq2_0_f16")]] kernel mul_mm_id kernel_mul_mm_id; +template [[host_name("kernel_mul_mm_id_ptq1_0_f16")]] kernel mul_mm_id kernel_mul_mm_id; template [[host_name("kernel_mul_mm_id_q4_0_f16")]] kernel mul_mm_id kernel_mul_mm_id; template [[host_name("kernel_mul_mm_id_q4_1_f16")]] kernel mul_mm_id kernel_mul_mm_id; template [[host_name("kernel_mul_mm_id_q5_0_f16")]] kernel mul_mm_id kernel_mul_mm_id; diff --git a/ggml/src/ggml-metal/kernels/mul_mv.metal b/ggml/src/ggml-metal/kernels/mul_mv.metal index bd4482ecf205..222902cc7e0b 100644 --- a/ggml/src/ggml-metal/kernels/mul_mv.metal +++ b/ggml/src/ggml-metal/kernels/mul_mv.metal @@ -510,6 +510,160 @@ kernel void kernel_mul_mv_q2_0_f32( kernel_mul_mv_q2_0_f32_impl(args, src0, src1, dst, nullptr, tgpig, tiisg, sgitg); } +// Byte-owning dot for PTQ1_0. The previous mapping gave each thread 16 contiguous +// weights, which in base-3 layout means every thread reloads the same qs bytes for a +// different trit index -- five times the byte traffic for the same 24 bytes. Here a +// thread instead owns whole bytes and consumes all five trits of each one, so a block's +// 26 bytes are loaded exactly once across the eight threads that cover it. The cost is +// that the y reads become strided rather than contiguous, which is the cheaper side: +// weight traffic is what decode is bound by. +// +// Byte ownership for thread it in 0..7, matching the CPU codec's element order: +// qs[2*it], qs[2*it+1] -> elements n*16 + m for n in 0..4 +// qs[16 + it] -> elements 80 + n*8 + it for n in 0..4 +// qh[it] (it < 2 only) -> elements 120 + n*2 + it for n in 0..3 +// Dot against y already staged in registers by the caller and reused across all nr0 +// rows. Trits come out of the byte by the base-3 remainder recurrence +// t = (v*3) >> 8, v = (v*3) & 0xFF, v starting at the byte +// which is two integer ops per trit with no table and no pow3 lookup. A 256-entry +// lookup was measurably worse here: it adds a dependent load per byte, and decode on +// this kernel is instruction-bound, not bandwidth-bound. Verified against the +// reference decode for all 256 bytes and all five positions. +// Accumulates the raw 0/1/2 trit and subtracts the staged y sum once, so there is no +// per-element offset correction. +inline float ptq1_0_dot_reg(device const block_ptq1_0 * qb, thread const float * yl, float sumy, short it) { + float acc = 0.f; + short c = 0; + + FOR_UNROLL (short k = 0; k < 2; ++k) { + ushort v = qb->qs[2*it + k]; + FOR_UNROLL (short n = 0; n < 5; ++n) { + const ushort w = v * 3; + acc += (float) (w >> 8) * yl[c++]; + v = w & 0xFF; + } + } + + { + ushort v = qb->qs[16 + it]; + FOR_UNROLL (short n = 0; n < 5; ++n) { + const ushort w = v * 3; + acc += (float) (w >> 8) * yl[c++]; + v = w & 0xFF; + } + } + + // qh holds 8 elements; give every thread exactly one so all eight do 16 elements. + // Element 120+it sits at trit it>>1 of byte qh[it&1], so step the recurrence to it. + { + ushort v = qb->qh[it & 1]; + const short n = it >> 1; + for (short i = 0; i < n; ++i) { + v = (v * 3) & 0xFF; + } + acc += (float) ((v * 3) >> 8) * yl[c++]; + } + + return (acc - sumy) * (float) qb->d; +} + +template +void kernel_mul_mv_ptq1_0_f32_impl( + args_t args, + device const char * src0, + device const char * src1, + device char * dst, + threadgroup char * shmem, + uint3 tgpig, + ushort tiisg, + ushort sgitg) { + const short NSG = FC_mul_mv_nsg; + + const int nb = args.ne00/QK_PTQ1_0; + + const int r0 = tgpig.x; + const int r1 = tgpig.y; + const int im = tgpig.z; + + const int first_row = (r0 * NSG + sgitg) * nr0; + + const uint i12 = im%FC_mul_mv_ne12; + const uint i13 = im/FC_mul_mv_ne12; + + const uint64_t offset1 = r1*args.nb11 + (i12)*args.nb12 + (i13)*args.nb13; + + device const float * y = (device const float *) (src1 + offset1); + + device const block_ptq1_0 * ax[nr0]; + for (int row = 0; row < nr0; ++row) { + const uint64_t offset0 = (first_row + row)*args.nb01 + (i12/FC_mul_mv_r2)*args.nb02 + (i13/FC_mul_mv_r3)*args.nb03; + ax[row] = (device const block_ptq1_0 *) ((device char *) src0 + offset0); + } + + float yl[16]; + float sumf[nr0] = {0.f}; + + // eight threads cover one 128-weight block; each owns whole bytes, not a + // contiguous element span, so the block's bytes are read once in total + const short ix = (tiisg/8); + const short it = (tiisg%8); + + device const float * yb = y + ix*QK_PTQ1_0; + + for (int ib = ix; ib < nb; ib += N_SIMDWIDTH/8) { + // stage this thread's y values once, then reuse them for every row + float sumy = 0.f; + short c = 0; + + FOR_UNROLL (short k = 0; k < 2; ++k) { + const short m = 2*it + k; + FOR_UNROLL (short n = 0; n < 5; ++n) { + const float v = yb[n*16 + m]; + yl[c++] = v; + sumy += v; + } + } + FOR_UNROLL (short n = 0; n < 5; ++n) { + const float v = yb[80 + n*8 + it]; + yl[c++] = v; + sumy += v; + } + { + const float v = yb[120 + it]; + yl[c++] = v; + sumy += v; + } + + FOR_UNROLL (short row = 0; row < nr0; row++) { + sumf[row] += ptq1_0_dot_reg(ax[row] + ib, yl, sumy, it); + } + + yb += QK_PTQ1_0 * (N_SIMDWIDTH/8); + } + + device float * dst_f32 = (device float *) dst + (uint64_t)im*args.ne0*args.ne1 + (uint64_t)r1*args.ne0; + + for (int row = 0; row < nr0; ++row) { + const float tot = simd_sum(sumf[row]); + + if (tiisg == 0 && first_row + row < args.ne01) { + dst_f32[first_row + row] = tot; + } + } +} + +[[host_name("kernel_mul_mv_ptq1_0_f32")]] +kernel void kernel_mul_mv_ptq1_0_f32( + constant ggml_metal_kargs_mul_mv & args, + device const char * src0, + device const char * src1, + device char * dst, + uint3 tgpig[[threadgroup_position_in_grid]], + ushort tiisg[[thread_index_in_simdgroup]], + ushort sgitg[[simdgroup_index_in_threadgroup]]) { + kernel_mul_mv_ptq1_0_f32_impl(args, src0, src1, dst, nullptr, tgpig, tiisg, sgitg); +} + template void kernel_mul_mv_pq2_0_f32_impl( args_t args, @@ -995,6 +1149,10 @@ template [[host_name("kernel_mul_mv_ext_q2_0_f32_r1_5")]] kernel mul_mv_ext_q4 template [[host_name("kernel_mul_mv_ext_pq2_0_f32_r1_2")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<2, block_pq2_0, 128, dequantize_pq2_0_t4>; template [[host_name("kernel_mul_mv_ext_pq2_0_f32_r1_3")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<3, block_pq2_0, 128, dequantize_pq2_0_t4>; template [[host_name("kernel_mul_mv_ext_pq2_0_f32_r1_4")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<4, block_pq2_0, 128, dequantize_pq2_0_t4>; +template [[host_name("kernel_mul_mv_ext_ptq1_0_f32_r1_2")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<2, block_ptq1_0, 128, dequantize_ptq1_0_t4>; +template [[host_name("kernel_mul_mv_ext_ptq1_0_f32_r1_3")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<3, block_ptq1_0, 128, dequantize_ptq1_0_t4>; +template [[host_name("kernel_mul_mv_ext_ptq1_0_f32_r1_4")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<4, block_ptq1_0, 128, dequantize_ptq1_0_t4>; +template [[host_name("kernel_mul_mv_ext_ptq1_0_f32_r1_5")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<5, block_ptq1_0, 128, dequantize_ptq1_0_t4>; template [[host_name("kernel_mul_mv_ext_pq2_0_f32_r1_5")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<5, block_pq2_0, 128, dequantize_pq2_0_t4>; template [[host_name("kernel_mul_mv_ext_q4_0_f32_r1_2")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<2, block_q4_0, 32, dequantize_q4_0_t4>; diff --git a/ggml/src/ggml-metal/kernels/quantize.metal b/ggml/src/ggml-metal/kernels/quantize.metal index c3ac9eccdbfd..a49d8ad7e234 100644 --- a/ggml/src/ggml-metal/kernels/quantize.metal +++ b/ggml/src/ggml-metal/kernels/quantize.metal @@ -284,6 +284,7 @@ typedef decltype(kernel_get_rows_q) get_rows_q_t template [[host_name("kernel_get_rows_q1_0")]] kernel get_rows_q_t kernel_get_rows_q; template [[host_name("kernel_get_rows_q2_0")]] kernel get_rows_q_t kernel_get_rows_q; template [[host_name("kernel_get_rows_pq2_0")]] kernel get_rows_q_t kernel_get_rows_q; +template [[host_name("kernel_get_rows_ptq1_0")]] kernel get_rows_q_t kernel_get_rows_q; template [[host_name("kernel_get_rows_q4_0")]] kernel get_rows_q_t kernel_get_rows_q; template [[host_name("kernel_get_rows_q4_1")]] kernel get_rows_q_t kernel_get_rows_q; template [[host_name("kernel_get_rows_q5_0")]] kernel get_rows_q_t kernel_get_rows_q; diff --git a/ggml/src/ggml-quants.c b/ggml/src/ggml-quants.c index a5f8f9c352f3..bfb3892a599c 100644 --- a/ggml/src/ggml-quants.c +++ b/ggml/src/ggml-quants.c @@ -2195,6 +2195,107 @@ size_t quantize_pq2_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, return nrow * row_size; } +// ====================== PTQ1_0 (Prism ternary, group 128) ====================== +// Base-3 trit packing identical to upstream TQ1_0, but at block 128 so one fp16 +// scale covers 128 weights. qs is 24 bytes, which TQ1_0's fixed 32-then-16 byte +// staging cannot cover, so the stages are generalised to 32/16/8; at TQ1_0's +// 48-byte qs this reduces to exactly its original 32-then-16 behaviour. +static const size_t ptq1_0_stages[3] = {32, 16, 8}; + +void quantize_row_ptq1_0_ref(const float * GGML_RESTRICT x, block_ptq1_0 * GGML_RESTRICT y, int64_t k) { + assert(k % QK_PTQ1_0 == 0); + const int64_t nb = k / QK_PTQ1_0; + + for (int64_t i = 0; i < nb; i++) { + float amax = 0.0f; + for (int j = 0; j < QK_PTQ1_0; j++) { + amax = MAX(amax, fabsf(x[j])); + } + + const float d = amax; + const float id = d ? 1.0f/d : 0.0f; + + y[i].d = GGML_FP32_TO_FP16(d); + + size_t j = 0; + for (size_t s = 0; s < 3; ++s) { + const size_t c = ptq1_0_stages[s]; + for (; j + c <= sizeof(y->qs); j += c) { + for (size_t m = 0; m < c; ++m) { + uint8_t q = 0; + for (size_t n = 0; n < 5; ++n) { + int xi = lroundf(x[m + n*c] * id) + 1; // -1, 0, 1 -> 0, 1, 2 + q *= 3; + q += xi; + } + // ceiling division (243 == pow(3, 5)) + q = ((uint16_t)q * 256 + (243 - 1)) / 243; + y[i].qs[j + m] = q; + } + x += 5*c; + } + } + // 4 elements per byte + for (size_t h = 0; h < sizeof(y->qh); ++h) { + uint8_t q = 0; + for (size_t m = 0; m < 4; ++m) { + int xi = lroundf(x[h + m*sizeof(y->qh)] * id) + 1; + q *= 3; + q += xi; + } + // shift the first value to the most significant trit + q *= 3; + q = ((uint16_t)q * 256 + (243 - 1)) / 243; + y[i].qh[h] = q; + } + x += 4*sizeof(y->qh); + } +} + +void dequantize_row_ptq1_0(const block_ptq1_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { + assert(k % QK_PTQ1_0 == 0); + const int64_t nb = k / QK_PTQ1_0; + + const uint8_t pow3[6] = {1, 3, 9, 27, 81, 243}; + + for (int64_t i = 0; i < nb; ++i) { + const float d = GGML_FP16_TO_FP32(x[i].d); + + size_t j = 0; + for (size_t s = 0; s < 3; ++s) { + const size_t c = ptq1_0_stages[s]; + for (; j + c <= sizeof(x->qs); j += c) { + for (size_t n = 0; n < 5; ++n) { + for (size_t m = 0; m < c; ++m) { + uint8_t q = x[i].qs[j + m] * pow3[n]; + int16_t xi = ((uint16_t) q * 3) >> 8; + *y++ = (float) (xi - 1) * d; + } + } + } + } + for (size_t n = 0; n < 4; ++n) { + for (size_t h = 0; h < sizeof(x->qh); ++h) { + uint8_t q = x[i].qh[h] * pow3[n]; + int16_t xi = ((uint16_t) q * 3) >> 8; + *y++ = (float) (xi - 1) * d; + } + } + } +} + +size_t quantize_ptq1_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) { + (void)quant_weights; // ternary codes come from the weights themselves; an imatrix has no role + const size_t row_size = ggml_row_size(GGML_TYPE_PTQ1_0, n_per_row); + char * qrow = (char *)dst; + for (int64_t row = 0; row < nrow; ++row) { + quantize_row_ptq1_0_ref(src, (block_ptq1_0 *)qrow, n_per_row); + src += n_per_row; + qrow += row_size; + } + return nrow * row_size; +} + size_t quantize_q4_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) { if (!quant_weights) { quantize_row_q4_0_ref(src, dst, (int64_t)nrow*n_per_row); @@ -5611,6 +5712,10 @@ bool ggml_validate_row_data(enum ggml_type type, const void * data, size_t nbyte { VALIDATE_ROW_DATA_D_F16_IMPL(block_pq2_0, data, nb); } break; + case GGML_TYPE_PTQ1_0: + { + VALIDATE_ROW_DATA_D_F16_IMPL(block_ptq1_0, data, nb); + } break; case GGML_TYPE_Q4_0: { VALIDATE_ROW_DATA_D_F16_IMPL(block_q4_0, data, nb); diff --git a/ggml/src/ggml-quants.h b/ggml/src/ggml-quants.h index 8a8370f48ef2..c9e5b2fe3587 100644 --- a/ggml/src/ggml-quants.h +++ b/ggml/src/ggml-quants.h @@ -17,6 +17,7 @@ extern "C" { GGML_API void quantize_row_q1_0_ref(const float * GGML_RESTRICT x, block_q1_0 * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_q2_0_ref(const float * GGML_RESTRICT x, block_q2_0 * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_pq2_0_ref(const float * GGML_RESTRICT x, block_pq2_0 * GGML_RESTRICT y, int64_t k); +GGML_API void quantize_row_ptq1_0_ref(const float * GGML_RESTRICT x, block_ptq1_0 * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_q4_0_ref(const float * GGML_RESTRICT x, block_q4_0 * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_q4_1_ref(const float * GGML_RESTRICT x, block_q4_1 * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_q5_0_ref(const float * GGML_RESTRICT x, block_q5_0 * GGML_RESTRICT y, int64_t k); @@ -47,6 +48,7 @@ GGML_API void quantize_row_iq2_s_ref (const float * GGML_RESTRICT x, block_iq2_ GGML_API void dequantize_row_q1_0(const block_q1_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); GGML_API void dequantize_row_q2_0(const block_q2_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); GGML_API void dequantize_row_pq2_0(const block_pq2_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API void dequantize_row_ptq1_0(const block_ptq1_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); GGML_API void dequantize_row_q4_0(const block_q4_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); GGML_API void dequantize_row_q4_1(const block_q4_1 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); GGML_API void dequantize_row_q5_0(const block_q5_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); @@ -99,6 +101,7 @@ GGML_API size_t quantize_q6_K(const float * GGML_RESTRICT src, void * GGML_RESTR GGML_API size_t quantize_q1_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); GGML_API size_t quantize_q2_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); GGML_API size_t quantize_pq2_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); +GGML_API size_t quantize_ptq1_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); GGML_API size_t quantize_q4_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); GGML_API size_t quantize_q4_1(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); GGML_API size_t quantize_q5_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp index 832e344ccffd..e30fb0d63329 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp +++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp @@ -4634,6 +4634,9 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { } #endif CREATE_MM2(pipeline_dequant_mul_mat_mat_f16[GGML_TYPE_Q1_0], matmul_q1_0_f16, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3) + // PTQ1_0 intentionally has no coopmat2 pipelines: dequant_funcs_cm2.glsl carries no + // PTQ1_0 decoder, so vulkan-shaders-gen skips cm2 generation for it and these symbols + // do not exist. It uses the scalar and coopmat1 matmul paths instead. CREATE_MM2(pipeline_dequant_mul_mat_mat_f16[GGML_TYPE_Q2_0], matmul_q2_0_f16, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3) CREATE_MM2(pipeline_dequant_mul_mat_mat_f16[GGML_TYPE_Q4_0], matmul_q4_0_f16, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3) CREATE_MM2(pipeline_dequant_mul_mat_mat_f16[GGML_TYPE_Q4_1], matmul_q4_1_f16, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3) @@ -4747,6 +4750,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { #endif CREATE_MM2(GGML_TYPE_Q1_0, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q1_0], matmul_q1_0_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, ); + CREATE_MM2(GGML_TYPE_PTQ1_0, pipeline_dequant_mul_mat_mat[GGML_TYPE_PTQ1_0], matmul_ptq1_0_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, ); CREATE_MM2(GGML_TYPE_Q2_0, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q2_0], matmul_q2_0_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, ); CREATE_MM2(GGML_TYPE_Q4_0, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q4_0], matmul_q4_0_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, ); CREATE_MM2(GGML_TYPE_Q4_1, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q4_1], matmul_q4_1_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, ); @@ -4793,6 +4797,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { #endif CREATE_MM2(GGML_TYPE_Q1_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q1_0], matmul_id_subgroup_q1_0_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id); + CREATE_MM2(GGML_TYPE_PTQ1_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_PTQ1_0], matmul_id_subgroup_ptq1_0_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id); CREATE_MM2(GGML_TYPE_Q2_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q2_0], matmul_id_subgroup_q2_0_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id); CREATE_MM2(GGML_TYPE_Q4_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q4_0], matmul_id_subgroup_q4_0_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id); CREATE_MM2(GGML_TYPE_Q4_1, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q4_1], matmul_id_subgroup_q4_1_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id); @@ -4884,6 +4889,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { CREATE_MM_NODOT2(GGML_TYPE_BF16, pipeline_matmul_bf16, matmul_bf16, , wg_denoms, warptile, vk_mat_mat_push_constants, 3, , 0); CREATE_MM2(GGML_TYPE_Q1_0, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q1_0], matmul_q1_0_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0); + CREATE_MM2(GGML_TYPE_PTQ1_0, pipeline_dequant_mul_mat_mat[GGML_TYPE_PTQ1_0], matmul_ptq1_0_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0); CREATE_MM2(GGML_TYPE_Q2_0, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q2_0], matmul_q2_0_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0); CREATE_MM2(GGML_TYPE_Q4_0, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q4_0], matmul_q4_0_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0); CREATE_MM2(GGML_TYPE_Q4_1, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q4_1], matmul_q4_1_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0); @@ -4933,6 +4939,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { CREATE_MM2(GGML_TYPE_F16, pipeline_matmul_id_f16_f32, matmul_id_subgroup_f16_f32, wg_denoms, warptile_id, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size_16); CREATE_MM_NODOT2(GGML_TYPE_BF16, pipeline_matmul_id_bf16, matmul_id_subgroup_bf16, , wg_denoms, warptile_id, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size_16); CREATE_MM2(GGML_TYPE_Q1_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q1_0], matmul_id_subgroup_q1_0_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size); + CREATE_MM2(GGML_TYPE_PTQ1_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_PTQ1_0], matmul_id_subgroup_ptq1_0_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size); CREATE_MM2(GGML_TYPE_Q2_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q2_0], matmul_id_subgroup_q2_0_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size); CREATE_MM2(GGML_TYPE_Q4_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q4_0], matmul_id_subgroup_q4_0_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size); CREATE_MM2(GGML_TYPE_Q4_1, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q4_1], matmul_id_subgroup_q4_1_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size); @@ -4981,6 +4988,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { CREATE_MM2(GGML_TYPE_F16, pipeline_matmul_id_f16_f32, matmul_id_f16_f32, wg_denoms, warptile, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0); CREATE_MM_NODOT2(GGML_TYPE_BF16, pipeline_matmul_id_bf16, matmul_id_bf16, , wg_denoms, warptile, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0); CREATE_MM2(GGML_TYPE_Q1_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q1_0], matmul_id_q1_0_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0); + CREATE_MM2(GGML_TYPE_PTQ1_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_PTQ1_0], matmul_id_ptq1_0_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0); CREATE_MM2(GGML_TYPE_Q2_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q2_0], matmul_id_q2_0_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0); CREATE_MM2(GGML_TYPE_Q4_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q4_0], matmul_id_q4_0_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0); CREATE_MM2(GGML_TYPE_Q4_1, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q4_1], matmul_id_q4_1_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0); @@ -5060,6 +5068,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { CREATE_MM(GGML_TYPE_BF16, pipeline_matmul_bf16, matmul_bf16, , wg_denoms, warptile, vk_mat_mat_push_constants, 3, , 0); CREATE_MM(GGML_TYPE_Q1_0, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q1_0].f32acc, matmul_q1_0_f32, , mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0); + CREATE_MM(GGML_TYPE_PTQ1_0, pipeline_dequant_mul_mat_mat[GGML_TYPE_PTQ1_0].f32acc, matmul_ptq1_0_f32, , mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0); CREATE_MM(GGML_TYPE_Q2_0, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q2_0].f32acc, matmul_q2_0_f32, , mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0); CREATE_MM(GGML_TYPE_Q4_0, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q4_0].f32acc, matmul_q4_0_f32, , mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0); CREATE_MM(GGML_TYPE_Q4_1, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q4_1].f32acc, matmul_q4_1_f32, , mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0); @@ -5109,6 +5118,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { CREATE_MM(GGML_TYPE_BF16, pipeline_matmul_id_bf16, matmul_id_subgroup_bf16, , wg_denoms, warptile_id, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size_16); CREATE_MM(GGML_TYPE_Q1_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q1_0].f32acc, matmul_id_subgroup_q1_0_f32, , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size); + CREATE_MM(GGML_TYPE_PTQ1_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_PTQ1_0].f32acc, matmul_id_subgroup_ptq1_0_f32, , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size); CREATE_MM(GGML_TYPE_Q2_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q2_0].f32acc, matmul_id_subgroup_q2_0_f32, , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size); CREATE_MM(GGML_TYPE_Q4_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q4_0].f32acc, matmul_id_subgroup_q4_0_f32, , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size); CREATE_MM(GGML_TYPE_Q4_1, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q4_1].f32acc, matmul_id_subgroup_q4_1_f32, , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size); @@ -5139,6 +5149,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { CREATE_MM(GGML_TYPE_BF16, pipeline_matmul_id_bf16, matmul_id_bf16, , wg_denoms, warptile, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0); CREATE_MM(GGML_TYPE_Q1_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q1_0].f32acc, matmul_id_q1_0_f32, , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0); + CREATE_MM(GGML_TYPE_PTQ1_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_PTQ1_0].f32acc, matmul_id_ptq1_0_f32, , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0); CREATE_MM(GGML_TYPE_Q2_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q2_0].f32acc, matmul_id_q2_0_f32, , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0); CREATE_MM(GGML_TYPE_Q4_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q4_0].f32acc, matmul_id_q4_0_f32, , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0); CREATE_MM(GGML_TYPE_Q4_1, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q4_1].f32acc, matmul_id_q4_1_f32, , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0); @@ -5243,6 +5254,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f32_f32[w][GGML_TYPE_F16 ][i], "mul_mat_vec_f16_f32_f32", arr_dmmv_f16_f32_f32_len[reduc], arr_dmmv_f16_f32_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2, 1, 1}, {wg_size_subgroup, 2, i+1}, 1, false, use_subgroups, force_subgroup_size); ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f32_f32[w][GGML_TYPE_BF16][i], "mul_mat_vec_bf16_f32_f32", arr_dmmv_bf16_f32_f32_len[reduc], arr_dmmv_bf16_f32_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2, 1, 1}, {wg_size_subgroup, 2, i+1}, 1, false, use_subgroups, force_subgroup_size); ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f32_f32[w][GGML_TYPE_Q1_0][i], "mul_mat_vec_q1_0_f32_f32", arr_dmmv_q1_0_f32_f32_len[reduc], arr_dmmv_q1_0_f32_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq, i+1}, 1, true, use_subgroups, force_subgroup_size); + ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f32_f32[w][GGML_TYPE_PTQ1_0][i], "mul_mat_vec_ptq1_0_f32_f32", arr_dmmv_ptq1_0_f32_f32_len[reduc], arr_dmmv_ptq1_0_f32_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq, i+1}, 1, true, use_subgroups, force_subgroup_size); ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f32_f32[w][GGML_TYPE_Q2_0][i], "mul_mat_vec_q2_0_f32_f32", arr_dmmv_q2_0_f32_f32_len[reduc], arr_dmmv_q2_0_f32_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq, i+1}, 1, true, use_subgroups, force_subgroup_size); ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f32_f32[w][GGML_TYPE_Q4_0][i], "mul_mat_vec_q4_0_f32_f32", arr_dmmv_q4_0_f32_f32_len[reduc], arr_dmmv_q4_0_f32_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq, i+1}, 1, true, use_subgroups, force_subgroup_size); ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f32_f32[w][GGML_TYPE_Q4_1][i], "mul_mat_vec_q4_1_f32_f32", arr_dmmv_q4_1_f32_f32_len[reduc], arr_dmmv_q4_1_f32_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq, i+1}, 1, true, use_subgroups, force_subgroup_size); @@ -5271,6 +5283,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f16_f32[w][GGML_TYPE_F16 ][i], "mul_mat_vec_f16_f16_f32", arr_dmmv_f16_f16_f32_len[reduc], arr_dmmv_f16_f16_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2, 1, 1}, {wg_size_subgroup, 2, i+1}, 1, false, use_subgroups, force_subgroup_size); ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f16_f32[w][GGML_TYPE_BF16][i], "mul_mat_vec_bf16_f16_f32", arr_dmmv_bf16_f16_f32_len[reduc], arr_dmmv_bf16_f16_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2, 1, 1}, {wg_size_subgroup, 2, i+1}, 1, false, use_subgroups, force_subgroup_size); ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f16_f32[w][GGML_TYPE_Q1_0][i], "mul_mat_vec_q1_0_f16_f32", arr_dmmv_q1_0_f16_f32_len[reduc], arr_dmmv_q1_0_f16_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq, i+1}, 1, true, use_subgroups, force_subgroup_size); + ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f16_f32[w][GGML_TYPE_PTQ1_0][i], "mul_mat_vec_ptq1_0_f16_f32", arr_dmmv_ptq1_0_f16_f32_len[reduc], arr_dmmv_ptq1_0_f16_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq, i+1}, 1, true, use_subgroups, force_subgroup_size); ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f16_f32[w][GGML_TYPE_Q2_0][i], "mul_mat_vec_q2_0_f16_f32", arr_dmmv_q2_0_f16_f32_len[reduc], arr_dmmv_q2_0_f16_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq, i+1}, 1, true, use_subgroups, force_subgroup_size); ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f16_f32[w][GGML_TYPE_Q4_0][i], "mul_mat_vec_q4_0_f16_f32", arr_dmmv_q4_0_f16_f32_len[reduc], arr_dmmv_q4_0_f16_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq, i+1}, 1, true, use_subgroups, force_subgroup_size); ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f16_f32[w][GGML_TYPE_Q4_1][i], "mul_mat_vec_q4_1_f16_f32", arr_dmmv_q4_1_f16_f32_len[reduc], arr_dmmv_q4_1_f16_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq, i+1}, 1, true, use_subgroups, force_subgroup_size); @@ -5326,6 +5339,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_id_f32[w][GGML_TYPE_F16 ], "mul_mat_vec_id_f16_f32", arr_dmmv_id_f16_f32_f32_len[reduc], arr_dmmv_id_f16_f32_f32_data[reduc], "main", mul_mat_vec_id_num_bindings, sizeof(vk_mat_vec_id_push_constants), {2, 1, 1}, {wg_size_subgroup, 2}, 1, false, use_subgroups, force_subgroup_size); ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_id_f32[w][GGML_TYPE_BF16], "mul_mat_vec_id_bf16_f32", arr_dmmv_id_bf16_f32_f32_len[reduc], arr_dmmv_id_bf16_f32_f32_data[reduc], "main", mul_mat_vec_id_num_bindings, sizeof(vk_mat_vec_id_push_constants), {2, 1, 1}, {wg_size_subgroup, 2}, 1, false, use_subgroups, force_subgroup_size); ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_id_f32[w][GGML_TYPE_Q1_0], "mul_mat_vec_id_q1_0_f32", arr_dmmv_id_q1_0_f32_f32_len[reduc], arr_dmmv_id_q1_0_f32_f32_data[reduc], "main", mul_mat_vec_id_num_bindings, sizeof(vk_mat_vec_id_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq}, 1, true, use_subgroups, force_subgroup_size); + ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_id_f32[w][GGML_TYPE_PTQ1_0], "mul_mat_vec_id_ptq1_0_f32", arr_dmmv_id_ptq1_0_f32_f32_len[reduc], arr_dmmv_id_ptq1_0_f32_f32_data[reduc], "main", mul_mat_vec_id_num_bindings, sizeof(vk_mat_vec_id_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq}, 1, true, use_subgroups, force_subgroup_size); ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_id_f32[w][GGML_TYPE_Q2_0], "mul_mat_vec_id_q2_0_f32", arr_dmmv_id_q2_0_f32_f32_len[reduc], arr_dmmv_id_q2_0_f32_f32_data[reduc], "main", mul_mat_vec_id_num_bindings, sizeof(vk_mat_vec_id_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq}, 1, true, use_subgroups, force_subgroup_size); ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_id_f32[w][GGML_TYPE_Q4_0], "mul_mat_vec_id_q4_0_f32", arr_dmmv_id_q4_0_f32_f32_len[reduc], arr_dmmv_id_q4_0_f32_f32_data[reduc], "main", mul_mat_vec_id_num_bindings, sizeof(vk_mat_vec_id_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq}, 1, true, use_subgroups, force_subgroup_size); ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_id_f32[w][GGML_TYPE_Q4_1], "mul_mat_vec_id_q4_1_f32", arr_dmmv_id_q4_1_f32_f32_len[reduc], arr_dmmv_id_q4_1_f32_f32_data[reduc], "main", mul_mat_vec_id_num_bindings, sizeof(vk_mat_vec_id_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq}, 1, true, use_subgroups, force_subgroup_size); @@ -5388,6 +5402,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { // dequant shaders ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_F32 ], "f32_to_f16", dequant_f32_len, dequant_f32_data, "main", 2, 5 * sizeof(uint32_t), {256 * 16, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_Q1_0], "dequant_q1_0", dequant_q1_0_len, dequant_q1_0_data, "main", 2, 5 * sizeof(uint32_t), {256 * 8, 1, 1}, {}, 1); + ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_PTQ1_0], "dequant_ptq1_0", dequant_ptq1_0_len, dequant_ptq1_0_data, "main", 2, 5 * sizeof(uint32_t), {256 * 8, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_Q2_0], "dequant_q2_0", dequant_q2_0_len, dequant_q2_0_data, "main", 2, 5 * sizeof(uint32_t), {256 * 4, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_Q4_0], "dequant_q4_0", dequant_q4_0_len, dequant_q4_0_data, "main", 2, 5 * sizeof(uint32_t), {256 * 16, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_Q4_1], "dequant_q4_1", dequant_q4_1_len, dequant_q4_1_data, "main", 2, 5 * sizeof(uint32_t), {256 * 16, 1, 1}, {}, 1); @@ -5418,6 +5433,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { ggml_vk_create_pipeline(device, device->pipeline_get_rows[GGML_TYPE_F16 ], "get_rows_f16", get_rows_f16_len, get_rows_f16_data, "main", 3, sizeof(vk_op_binary_push_constants), { 512, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_get_rows[GGML_TYPE_BF16], "get_rows_bf16", get_rows_bf16_len, get_rows_bf16_data, "main", 3, sizeof(vk_op_binary_push_constants), { 512, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_get_rows[GGML_TYPE_Q1_0], "get_rows_q1_0", get_rows_q1_0_len, get_rows_q1_0_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1); + ggml_vk_create_pipeline(device, device->pipeline_get_rows[GGML_TYPE_PTQ1_0], "get_rows_ptq1_0", get_rows_ptq1_0_len, get_rows_ptq1_0_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_get_rows[GGML_TYPE_Q2_0], "get_rows_q2_0", get_rows_q2_0_len, get_rows_q2_0_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_get_rows[GGML_TYPE_Q4_0], "get_rows_q4_0", get_rows_q4_0_len, get_rows_q4_0_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_get_rows[GGML_TYPE_Q4_1], "get_rows_q4_1", get_rows_q4_1_len, get_rows_q4_1_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1); @@ -5447,6 +5463,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { ggml_vk_create_pipeline(device, device->pipeline_get_rows_f32[GGML_TYPE_F16 ], "get_rows_f16_f32", get_rows_f16_f32_len, get_rows_f16_f32_data, "main", 3, sizeof(vk_op_binary_push_constants), { 512, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_get_rows_f32[GGML_TYPE_BF16], "get_rows_bf16_f32", get_rows_bf16_f32_len, get_rows_bf16_f32_data, "main", 3, sizeof(vk_op_binary_push_constants), { 512, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_get_rows_f32[GGML_TYPE_Q1_0], "get_rows_q1_0_f32", get_rows_q1_0_f32_len, get_rows_q1_0_f32_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1); + ggml_vk_create_pipeline(device, device->pipeline_get_rows_f32[GGML_TYPE_PTQ1_0], "get_rows_ptq1_0_f32", get_rows_ptq1_0_f32_len, get_rows_ptq1_0_f32_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_get_rows_f32[GGML_TYPE_Q2_0], "get_rows_q2_0_f32", get_rows_q2_0_f32_len, get_rows_q2_0_f32_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_get_rows_f32[GGML_TYPE_Q4_0], "get_rows_q4_0_f32", get_rows_q4_0_f32_len, get_rows_q4_0_f32_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_get_rows_f32[GGML_TYPE_Q4_1], "get_rows_q4_1_f32", get_rows_q4_1_f32_len, get_rows_q4_1_f32_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1); @@ -7656,6 +7673,7 @@ static vk_pipeline ggml_vk_get_to_fp16(ggml_backend_vk_context * ctx, ggml_type switch (type) { case GGML_TYPE_F32: case GGML_TYPE_Q1_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q2_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: @@ -7731,6 +7749,7 @@ static vk_matmul_pipeline ggml_vk_get_mul_mat_mat_pipeline(ggml_backend_vk_conte switch (src0_type) { case GGML_TYPE_Q1_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q2_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: @@ -7761,7 +7780,17 @@ static vk_matmul_pipeline ggml_vk_get_mul_mat_mat_pipeline(ggml_backend_vk_conte if (ctx->device->coopmat2) { assert(src1_type == GGML_TYPE_F16); - return prec == GGML_PREC_DEFAULT ? ctx->device->pipeline_dequant_mul_mat_mat_f16[src0_type].f16acc : ctx->device->pipeline_dequant_mul_mat_mat_f16[src0_type].f32acc; + vk_matmul_pipeline2 & mmp = ctx->device->pipeline_dequant_mul_mat_mat_f16[src0_type]; + vk_matmul_pipeline pipelines = prec == GGML_PREC_DEFAULT ? mmp.f16acc : mmp.f32acc; + + // A type with no coopmat2 decoder (PTQ1_0) leaves these slots empty. The slots are + // always allocated, so returning one would hand back a pipeline whose l/m/s are null; + // nullptr instead routes the caller to its dequant + f16 matmul fallback. + if (pipelines->is_empty()) { + return nullptr; + } + + return pipelines; } if (ctx->device->coopmat_support) { return (ctx->device->fp16 && ctx->device->coopmat_acc_f16_support && prec == GGML_PREC_DEFAULT) ? ctx->device->pipeline_dequant_mul_mat_mat[src0_type].f16acc : ctx->device->pipeline_dequant_mul_mat_mat[src0_type].f32acc; @@ -7801,6 +7830,7 @@ static vk_pipeline ggml_vk_get_dequantize_mul_mat_vec(ggml_backend_vk_context * case GGML_TYPE_F16: case GGML_TYPE_BF16: case GGML_TYPE_Q1_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q2_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: @@ -7895,6 +7925,7 @@ static vk_matmul_pipeline ggml_vk_get_mul_mat_mat_id_pipeline(ggml_backend_vk_co switch (src0_type) { case GGML_TYPE_Q1_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q2_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: @@ -7929,6 +7960,11 @@ static vk_matmul_pipeline ggml_vk_get_mul_mat_mat_id_pipeline(ggml_backend_vk_co bool support_fp16acc = !mmp.f16acc->is_empty(); bool support_fp32acc = !mmp.f32acc->is_empty(); + if (!support_fp16acc && !support_fp32acc) { + // No coopmat2 decoder for this type; let the caller dequantize to f16 instead of asserting. + return nullptr; + } + if (support_fp16acc && (prefer_fp16acc || !support_fp32acc)) { return mmp.f16acc; } else { @@ -7968,6 +8004,7 @@ static vk_pipeline ggml_vk_get_dequantize_mul_mat_vec_id(ggml_backend_vk_context case GGML_TYPE_F16: case GGML_TYPE_BF16: case GGML_TYPE_Q1_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q2_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: @@ -18148,6 +18185,7 @@ static bool ggml_backend_vk_device_supports_op(ggml_backend_dev_t dev, const ggm case GGML_TYPE_F16: case GGML_TYPE_BF16: case GGML_TYPE_Q1_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q2_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: @@ -18254,6 +18292,7 @@ static bool ggml_backend_vk_device_supports_op(ggml_backend_dev_t dev, const ggm case GGML_TYPE_F16: case GGML_TYPE_BF16: case GGML_TYPE_Q1_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q2_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/dequant_funcs.glsl b/ggml/src/ggml-vulkan/vulkan-shaders/dequant_funcs.glsl index 627932bd3547..88be1360403a 100644 --- a/ggml/src/ggml-vulkan/vulkan-shaders/dequant_funcs.glsl +++ b/ggml/src/ggml-vulkan/vulkan-shaders/dequant_funcs.glsl @@ -126,6 +126,20 @@ vec4 dequantize4(uint ib, uint iqs, uint a_offset) { } #endif +#if defined(DATA_A_PTQ1_0) +#include "ptq1_0.glsl" + +vec2 dequantize(uint ib, uint iqs, uint a_offset) { + return vec2(ptq1_0_trit(ib, a_offset, iqs), ptq1_0_trit(ib, a_offset, iqs + 1u)); +} +vec4 dequantize4(uint ib, uint iqs, uint a_offset) { + return vec4(ptq1_0_trit(ib, a_offset, iqs), + ptq1_0_trit(ib, a_offset, iqs + 1u), + ptq1_0_trit(ib, a_offset, iqs + 2u), + ptq1_0_trit(ib, a_offset, iqs + 3u)); +} +#endif + #if defined(DATA_A_Q1_0) vec2 dequantize(uint ib, uint iqs, uint a_offset) { const uint bits = uint(data_a[a_offset + ib].qs[iqs / 8u]) >> (iqs % 8u); @@ -564,6 +578,13 @@ vec2 get_dm(uint ib, uint a_offset) { } #endif +#if defined(DATA_A_PTQ1_0) +vec2 get_dm(uint ib, uint a_offset) { + const float d = float(data_a[a_offset + ib].d); + return vec2(d, 0); +} +#endif + #if defined(DATA_A_Q1_0) vec2 get_dm(uint ib, uint a_offset) { const float d = float(data_a[a_offset + ib].d); diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/dequant_ptq1_0.comp b/ggml/src/ggml-vulkan/vulkan-shaders/dequant_ptq1_0.comp new file mode 100644 index 000000000000..53baf41f65b9 --- /dev/null +++ b/ggml/src/ggml-vulkan/vulkan-shaders/dequant_ptq1_0.comp @@ -0,0 +1,54 @@ +#version 450 + +#include "dequant_head.glsl" + +layout(local_size_x = 256, local_size_y = 1, local_size_z = 1) in; + +layout (binding = 0) readonly buffer A {block_ptq1_0 data_a[];}; +layout (binding = 1) writeonly buffer D {D_TYPE data_b[];}; + +// One invocation per 8 elements: 16 invocations cover a 128-weight block. Element +// order follows the CPU codec (16-byte qs chunk, 8-byte qs chunk, then qh), so the +// index maths lives in one place rather than being reproduced per shader. +float ptq1_0_trit_l(uint ib, uint e) { + uint b; + uint n; + if (e < 80u) { + b = uint(data_a[ib].qs[e & 15u]); + n = e >> 4u; + } else if (e < 120u) { + const uint t = e - 80u; + b = uint(data_a[ib].qs[16u + (t & 7u)]); + n = t >> 3u; + } else { + const uint t = e - 120u; + b = uint(data_a[ib].qh[t & 1u]); + n = t >> 1u; + } + + uint v = b; + for (uint i = 0u; i < n; ++i) { + v = (v * 3u) & 0xFFu; + } + return float(int((v * 3u) >> 8u) - 1); +} + +void main() { + const uint i = gl_WorkGroupID.x * 4 + gl_LocalInvocationID.x / 64; + + const uint tid = gl_LocalInvocationID.x % 64; + const uint il = tid / 4; + const uint ir = tid % 4; + const uint ib = 4*i + ir; + if (ib >= p.nel / 128) { + return; + } + + const uint b_idx = 512*i + 128*ir + 8*il; + + const float d = float(data_a[ib].d); + + [[unroll]] for (uint l = 0; l < 8; ++l) { + data_b[b_idx + l] = D_TYPE(ptq1_0_trit_l(ib, 8*il + l) * d); + } +} diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/mul_mm.comp b/ggml/src/ggml-vulkan/vulkan-shaders/mul_mm.comp index 3df88044a5ee..9804c51a3f5a 100644 --- a/ggml/src/ggml-vulkan/vulkan-shaders/mul_mm.comp +++ b/ggml/src/ggml-vulkan/vulkan-shaders/mul_mm.comp @@ -136,6 +136,9 @@ shared FLOAT_TYPEV2 buf_b[BN * SHMEM_STRIDE]; shared ACC_TYPE coopmat_stage[TM * TN * NUM_WARPS]; #endif +#if defined(DATA_A_PTQ1_0) +#include "ptq1_0.glsl" +#endif #include "mul_mm_id_funcs.glsl" #include "mul_mm_funcs.glsl" diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/mul_mm_funcs.glsl b/ggml/src/ggml-vulkan/vulkan-shaders/mul_mm_funcs.glsl index 7d852dced8ab..3125a3fcfa93 100644 --- a/ggml/src/ggml-vulkan/vulkan-shaders/mul_mm_funcs.glsl +++ b/ggml/src/ggml-vulkan/vulkan-shaders/mul_mm_funcs.glsl @@ -152,6 +152,21 @@ void load_a_to_shmem(const uint pos_a, const uint row, const uint col, const uin const uint k_pair = row * LOAD_VEC_A / 2; store_a(col, k_pair, FLOAT_TYPEV2(v.xy)); store_a(col, k_pair + 1, FLOAT_TYPEV2(v.zw)); +#elif defined(DATA_A_PTQ1_0) + const uint idx = pos_a + col * p.stride_a / LOAD_VEC_A + row; + + const uint ib = idx / 16; + const uint grp = idx & 0xfu; // which 8-element group inside the block + const uint e0 = grp * 8u; + + const float d = float(data_a[ib].d); + + const uint k_pair = row * LOAD_VEC_A / 2; + [[unroll]] for (uint l = 0; l < 4; ++l) { + store_a(col, k_pair + l, FLOAT_TYPEV2( + ptq1_0_trit(ib, 0u, e0 + 2u*l) * d, + ptq1_0_trit(ib, 0u, e0 + 2u*l + 1u) * d)); + } #elif defined(DATA_A_Q1_0) const uint idx = pos_a + col * p.stride_a / LOAD_VEC_A + row; diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/ptq1_0.glsl b/ggml/src/ggml-vulkan/vulkan-shaders/ptq1_0.glsl new file mode 100644 index 000000000000..7eedd832b14f --- /dev/null +++ b/ggml/src/ggml-vulkan/vulkan-shaders/ptq1_0.glsl @@ -0,0 +1,38 @@ +#ifndef PTQ1_0_GLSL +#define PTQ1_0_GLSL + +// Shared PTQ1_0 trit accessor. Lives in its own header because two consumers need it +// from different include chains: dequant_funcs.glsl (mul_mat_vec, get_rows, +// copy_from_quant) and mul_mm.comp (via mul_mm_funcs.glsl), and mul_mm.comp does not +// include dequant_funcs.glsl. Duplicating it would leave two copies that must stay in +// step with the CPU codec in ggml-quants.c, where a divergence shows up as wrong +// matmul results rather than a build error. +// +// Element order is not positional: a 16-byte chunk of qs where byte j carries +// elements t*16+j, then an 8-byte chunk carrying 80 + t*8 + (j-16), then qh at four +// trits per byte carrying 120 + t*2 + h. Trits come out by the base-3 remainder +// recurrence t = (v*3)>>8, v = (v*3)&0xFF. +float ptq1_0_trit(uint ib, uint a_offset, uint e) { + uint b; + uint n; + if (e < 80u) { + b = uint(data_a[a_offset + ib].qs[e & 15u]); + n = e >> 4u; + } else if (e < 120u) { + const uint t = e - 80u; + b = uint(data_a[a_offset + ib].qs[16u + (t & 7u)]); + n = t >> 3u; + } else { + const uint t = e - 120u; + b = uint(data_a[a_offset + ib].qh[t & 1u]); + n = t >> 1u; + } + + uint v = b; + for (uint i = 0u; i < n; ++i) { + v = (v * 3u) & 0xFFu; + } + return float(int((v * 3u) >> 8u) - 1); +} + +#endif // PTQ1_0_GLSL diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/types.glsl b/ggml/src/ggml-vulkan/vulkan-shaders/types.glsl index adb1bb8b32b5..469aad51ce89 100644 --- a/ggml/src/ggml-vulkan/vulkan-shaders/types.glsl +++ b/ggml/src/ggml-vulkan/vulkan-shaders/types.glsl @@ -211,6 +211,27 @@ struct block_q1_0 #define A_TYPE block_q1_0 #endif +// PTQ1_0: ternary at group 128, base-3 packed five trits per byte. +// Field order mirrors block_ptq1_0 in ggml-common.h EXACTLY -- qs, then qh, then d. +// Unlike q1_0 the scale is LAST, and getting that wrong silently misindexes every +// block rather than failing loudly. +#define QUANT_K_PTQ1_0 128 +#define QUANT_R_PTQ1_0 1 + +struct block_ptq1_0 +{ + uint8_t qs[24]; + uint8_t qh[2]; + float16_t d; +}; + +#if defined(DATA_A_PTQ1_0) +#define QUANT_K QUANT_K_PTQ1_0 +#define QUANT_R QUANT_R_PTQ1_0 +#define QUANT_AUXF 1 +#define A_TYPE block_ptq1_0 +#endif + #define QUANT_K_Q2_0 64 #define QUANT_R_Q2_0 1 diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp b/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp index 770496b112c6..83af8af068fe 100644 --- a/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp +++ b/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp @@ -50,6 +50,7 @@ const std::vector type_names = { "f32", "f16", "q1_0", + "ptq1_0", "q2_0", "q4_0", "q4_1", @@ -583,7 +584,7 @@ void matmul_shaders(bool fp16, MatMulIdType matmul_id_type, bool coopmat, bool c for (const auto& tname : type_names) { std::string load_vec_quant = "2"; - if ((tname == "q1_0") || (tname == "q4_0") || (tname == "q4_1") || (tname == "q5_1") || (tname == "iq1_s") || (tname == "iq1_m") || (tname == "iq2_xxs") || (tname == "iq2_xs") || (tname == "iq2_s")) + if ((tname == "q1_0") || (tname == "ptq1_0") || (tname == "q4_0") || (tname == "q4_1") || (tname == "q5_1") || (tname == "iq1_s") || (tname == "iq1_m") || (tname == "iq2_xxs") || (tname == "iq2_xs") || (tname == "iq2_s")) load_vec_quant = "8"; else if ((tname == "q2_0") || (tname == "q5_0") || (tname == "q8_0") || (tname == "q2_k") || (tname == "q4_k") || (tname == "q5_k") || (tname == "iq3_xxs") || (tname == "iq3_s") || (tname == "iq4_xs") || (tname == "iq4_nl") || (tname == "mxfp4") || (tname == "nvfp4")) load_vec_quant = "4"; @@ -591,6 +592,13 @@ void matmul_shaders(bool fp16, MatMulIdType matmul_id_type, bool coopmat, bool c if (tname == "bf16") { continue; } + // PTQ1_0 has no coopmat2 decoder: dequant_funcs_cm2.glsl carries no PTQ1_0 entry, + // so emitting mul_mm_cm2 for it fails shader compilation and takes the whole + // Vulkan build down, not just this type. Skip it; it falls back to the scalar and + // coopmat1 matmul paths, which are the ones implemented and tested. + if (coopmat2 && tname == "ptq1_0") { + continue; + } std::string data_a_key = "DATA_A_" + to_uppercase(tname); // For aligned matmul loads diff --git a/ggml/src/ggml.c b/ggml/src/ggml.c index 6849721d8c60..2e7416070d04 100644 --- a/ggml/src/ggml.c +++ b/ggml/src/ggml.c @@ -697,6 +697,14 @@ static const struct ggml_type_traits type_traits[GGML_TYPE_COUNT] = { .to_float = (ggml_to_float_t) dequantize_row_pq2_0, .from_float_ref = (ggml_from_float_t) quantize_row_pq2_0_ref, }, + [GGML_TYPE_PTQ1_0] = { + .type_name = "ptq1_0", + .blck_size = QK_PTQ1_0, + .type_size = sizeof(block_ptq1_0), + .is_quantized = true, + .to_float = (ggml_to_float_t) dequantize_row_ptq1_0, + .from_float_ref = (ggml_from_float_t) quantize_row_ptq1_0_ref, + }, [GGML_TYPE_Q4_0] = { .type_name = "q4_0", .blck_size = QK4_0, @@ -1443,6 +1451,7 @@ enum ggml_type ggml_ftype_to_ggml_type(enum ggml_ftype ftype) { case GGML_FTYPE_MOSTLY_Q1_0: wtype = GGML_TYPE_Q1_0; break; case GGML_FTYPE_MOSTLY_Q2_0: wtype = GGML_TYPE_Q2_0; break; case GGML_FTYPE_MOSTLY_PQ2_0: wtype = GGML_TYPE_PQ2_0; break; + case GGML_FTYPE_MOSTLY_PTQ1_0: wtype = GGML_TYPE_PTQ1_0; break; case GGML_FTYPE_MOSTLY_Q5_0: wtype = GGML_TYPE_Q5_0; break; case GGML_FTYPE_MOSTLY_Q5_1: wtype = GGML_TYPE_Q5_1; break; case GGML_FTYPE_MOSTLY_Q8_0: wtype = GGML_TYPE_Q8_0; break; @@ -8050,6 +8059,7 @@ size_t ggml_quantize_chunk( case GGML_TYPE_Q1_0: result = quantize_q1_0 (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_Q2_0: result = quantize_q2_0 (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_PQ2_0: result = quantize_pq2_0(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; + case GGML_TYPE_PTQ1_0: result = quantize_ptq1_0(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_Q4_0: result = quantize_q4_0 (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_Q4_1: result = quantize_q4_1 (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_Q5_0: result = quantize_q5_0 (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; diff --git a/gguf-py/gguf/constants.py b/gguf-py/gguf/constants.py index 7777e7efa94d..d2284150ab6b 100644 --- a/gguf-py/gguf/constants.py +++ b/gguf-py/gguf/constants.py @@ -5384,7 +5384,8 @@ class GGMLQuantizationType(IntEnum): NVFP4 = 40 Q1_0 = 41 Q2_0 = 42 - PQ2_0 = 142 + PQ2_0 = 142 + PTQ1_0 = 143 class ExpertGatingFuncType(IntEnum): @@ -5441,7 +5442,8 @@ class LlamaFileType(IntEnum): MOSTLY_NVFP4 = 39 # except 1d tensors MOSTLY_Q1_0 = 40 # except 1d tensors MOSTLY_Q2_0 = 41 # except 1d tensors - MOSTLY_PQ2_0 = 128 # except 1d tensors + MOSTLY_PQ2_0 = 128 # except 1d tensors + MOSTLY_PTQ1_0 = 129 # except 1d tensors GUESSED = 1024 # not specified in the model file @@ -5578,7 +5580,8 @@ class VisionProjectorType: GGMLQuantizationType.NVFP4: (64, 4 + 32), GGMLQuantizationType.Q1_0: (128, 2 + 16), GGMLQuantizationType.Q2_0: (64, 2 + 16), - GGMLQuantizationType.PQ2_0: (128, 2 + 32), + GGMLQuantizationType.PQ2_0: (128, 2 + 32), + GGMLQuantizationType.PTQ1_0: (128, 2 + 24 + 2), } diff --git a/include/llama.h b/include/llama.h index e424d5c40100..d1aba602816a 100644 --- a/include/llama.h +++ b/include/llama.h @@ -158,6 +158,7 @@ extern "C" { LLAMA_FTYPE_MOSTLY_Q2_0 = 41, // except 1d tensors LLAMA_FTYPE_MOSTLY_PQ2_0 = 141, // except 1d tensors (Prism group-128 Q2_0; matches published PQ2_0 ggufs) LLAMA_FTYPE_MOSTLY_PQ2_0_LEGACY = 142, // pre-rename value for the same format, still found in published ggufs + LLAMA_FTYPE_MOSTLY_PTQ1_0 = 143, // except 1d tensors (Prism group-128 ternary, 1.75 bpw) LLAMA_FTYPE_GUESSED = 1024, // not specified in the model file }; diff --git a/src/llama-model-loader.cpp b/src/llama-model-loader.cpp index c43499aa234a..835d172d33f7 100644 --- a/src/llama-model-loader.cpp +++ b/src/llama-model-loader.cpp @@ -40,6 +40,7 @@ const char * llama_ftype_name(llama_ftype ftype) { case LLAMA_FTYPE_MOSTLY_Q1_0: name = LLAMA_FTYPE_PREFIX "Q1_0"; break; case LLAMA_FTYPE_MOSTLY_Q2_0: name = LLAMA_FTYPE_PREFIX "Q2_0"; break; case LLAMA_FTYPE_MOSTLY_PQ2_0: name = LLAMA_FTYPE_PREFIX "PQ2_0 - 2.13 bpw (group 128)"; break; + case LLAMA_FTYPE_MOSTLY_PTQ1_0: name = LLAMA_FTYPE_PREFIX "PTQ1_0 - 1.75 bpw ternary (group 128)"; break; // ggufs packed before the Q2_0_G128 -> PQ2_0 rename carry the old ftype value. // They load and compute correctly; name it so it does not report as unknown. case LLAMA_FTYPE_MOSTLY_PQ2_0_LEGACY: name = LLAMA_FTYPE_PREFIX "PQ2_0 - 2.13 bpw (group 128, legacy ftype)"; break; @@ -775,6 +776,7 @@ llama_model_loader::llama_model_loader( case GGML_TYPE_Q1_0: ftype = LLAMA_FTYPE_MOSTLY_Q1_0; break; case GGML_TYPE_Q2_0: ftype = LLAMA_FTYPE_MOSTLY_Q2_0; break; case GGML_TYPE_PQ2_0: ftype = LLAMA_FTYPE_MOSTLY_PQ2_0; break; + case GGML_TYPE_PTQ1_0: ftype = LLAMA_FTYPE_MOSTLY_PTQ1_0; break; default: { LLAMA_LOG_WARN("%s: unknown type %s\n", __func__, ggml_type_name(type_max)); diff --git a/src/llama-quant.cpp b/src/llama-quant.cpp index a06b44929c40..121e610abf5a 100644 --- a/src/llama-quant.cpp +++ b/src/llama-quant.cpp @@ -393,6 +393,7 @@ static ggml_type tensor_type_fallback(quantize_state_impl & qs, const ggml_tenso case GGML_TYPE_IQ3_S: // types on the right: block size 32 case GGML_TYPE_IQ4_XS: return_type = GGML_TYPE_IQ4_NL; break; case GGML_TYPE_Q2_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_PQ2_0: case GGML_TYPE_Q2_K: case GGML_TYPE_Q3_K: @@ -499,7 +500,7 @@ static ggml_type llama_tensor_get_type_impl(quantize_state_impl & qs, ggml_type else if (ftype == LLAMA_FTYPE_MOSTLY_IQ3_XXS) { new_type = GGML_TYPE_IQ3_S; } - else if (ftype == LLAMA_FTYPE_MOSTLY_TQ1_0 || ftype == LLAMA_FTYPE_MOSTLY_TQ2_0 || ftype == LLAMA_FTYPE_MOSTLY_Q2_0 || ftype == LLAMA_FTYPE_MOSTLY_PQ2_0) { + else if (ftype == LLAMA_FTYPE_MOSTLY_TQ1_0 || ftype == LLAMA_FTYPE_MOSTLY_TQ2_0 || ftype == LLAMA_FTYPE_MOSTLY_Q2_0 || ftype == LLAMA_FTYPE_MOSTLY_PQ2_0 || ftype == LLAMA_FTYPE_MOSTLY_PTQ1_0) { new_type = GGML_TYPE_Q4_K; } } @@ -821,6 +822,7 @@ ggml_type llama_ftype_get_default_type(llama_ftype ftype) { case LLAMA_FTYPE_MOSTLY_Q1_0: return GGML_TYPE_Q1_0; case LLAMA_FTYPE_MOSTLY_Q2_0: return GGML_TYPE_Q2_0; case LLAMA_FTYPE_MOSTLY_PQ2_0: return GGML_TYPE_PQ2_0; + case LLAMA_FTYPE_MOSTLY_PTQ1_0: return GGML_TYPE_PTQ1_0; case LLAMA_FTYPE_MOSTLY_MXFP4_MOE: return GGML_TYPE_MXFP4; diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index cca08f5dc5fd..16e3c74c3d9f 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -8338,7 +8338,7 @@ static const ggml_type all_types[] = { GGML_TYPE_Q8_0, GGML_TYPE_Q1_0, GGML_TYPE_Q2_0, - GGML_TYPE_PQ2_0, + GGML_TYPE_PQ2_0, GGML_TYPE_PTQ1_0, GGML_TYPE_MXFP4, GGML_TYPE_NVFP4, GGML_TYPE_Q2_K, GGML_TYPE_Q3_K, GGML_TYPE_Q4_K, GGML_TYPE_Q5_K, @@ -8355,7 +8355,7 @@ static const ggml_type base_types[] = { GGML_TYPE_Q8_0, // for I8MM tests GGML_TYPE_Q1_0, GGML_TYPE_Q2_0, - GGML_TYPE_PQ2_0, + GGML_TYPE_PQ2_0, GGML_TYPE_PTQ1_0, GGML_TYPE_Q4_0, GGML_TYPE_Q4_1, // for I8MM tests GGML_TYPE_Q4_K, @@ -8369,7 +8369,7 @@ static const ggml_type other_types[] = { GGML_TYPE_Q8_0, GGML_TYPE_Q1_0, GGML_TYPE_Q2_0, - GGML_TYPE_PQ2_0, + GGML_TYPE_PQ2_0, GGML_TYPE_PTQ1_0, GGML_TYPE_Q2_K, GGML_TYPE_Q3_K, GGML_TYPE_Q5_K, GGML_TYPE_Q6_K, diff --git a/tests/test-ptq1_0-cuda-dot.cpp b/tests/test-ptq1_0-cuda-dot.cpp new file mode 100644 index 000000000000..b94c279761f0 --- /dev/null +++ b/tests/test-ptq1_0-cuda-dot.cpp @@ -0,0 +1,132 @@ +// Verifies vec_dot_ptq1_0_q8_1 numerically on the host by stubbing the CUDA intrinsics. +// nvcc is unavailable here, so this cannot prove the kernel compiles, but it does prove +// the dot-product math and the element mapping agree with the CPU codec on real data. +#include +#include +#include +#include +#include + +// ---- CUDA intrinsic stubs ------------------------------------------------- +static inline int ggml_cuda_dp4a(int a, int b, int c) { // __dp4a + const int8_t* pa = (const int8_t*)&a; const int8_t* pb = (const int8_t*)&b; + return c + pa[0]*pb[0] + pa[1]*pb[1] + pa[2]*pb[2] + pa[3]*pb[3]; +} +struct half2_stub { float lo, hi; }; +static inline float __low2float(half2_stub h) { return h.lo; } + +#define QK8_1 32 +struct block_q8_1 { half2_stub ds; int8_t qs[QK8_1]; }; +static inline int get_int_b4(const int8_t* qs, int j) { // 4 bytes as an int + int v; memcpy(&v, qs + j*4, 4); return v; +} + +#define QK_PTQ1_0 128 +struct block_ptq1_0 { uint8_t qs[24]; uint8_t qh[2]; float d; }; + +// ---- transcribed from ggml-cuda/dequantize.cuh --------------------------- +static inline int ptq1_0_trit(const block_ptq1_0 * x, const int e) { + uint8_t b; int n; + if (e < 80) { b = x->qs[e & 15]; n = e >> 4; } + else if (e < 120) { const int t = e - 80; b = x->qs[16 + (t & 7)]; n = t >> 3; } + else { const int t = e - 120; b = x->qh[t & 1]; n = t >> 1; } + uint32_t v = b; + for (int i = 0; i < 4; ++i) if (i < n) v = (v * 3) & 0xFF; + return (int)((v * 3) >> 8) - 1; +} + +// ---- transcribed from ggml-cuda/vecdotq.cuh ------------------------------ +static inline float vec_dot_ptq1_0_q8_1(const void* vbq, const block_q8_1* bq8_1, + const int& kbx, const int& iqs) { + const block_ptq1_0 * bq = (const block_ptq1_0 *) vbq + kbx; + const block_q8_1 * bq8 = bq8_1 + iqs; + const int base = iqs * 32; + int sumi = 0; + for (int j = 0; j < 8; ++j) { + const int t0 = ptq1_0_trit(bq, base + j*4 + 0); + const int t1 = ptq1_0_trit(bq, base + j*4 + 1); + const int t2 = ptq1_0_trit(bq, base + j*4 + 2); + const int t3 = ptq1_0_trit(bq, base + j*4 + 3); + const int qx = (t0 & 0xFF) | ((t1 & 0xFF) << 8) | ((t2 & 0xFF) << 16) | ((t3 & 0xFF) << 24); + const int u = get_int_b4(bq8->qs, j); + sumi = ggml_cuda_dp4a(u, qx, sumi); + } + return (float) bq->d * __low2float(bq8->ds) * sumi; +} + +// ---- reference: dequantize the block, dequantize q8_1, dot in float ------ +static void ref_dequant(const block_ptq1_0* x, float* out) { + const uint8_t pow3[6]={1,3,9,27,81,243}; const size_t st[3]={32,16,8}; + int o=0; size_t j=0; + for (size_t s=0;s<3;++s){ const size_t c=st[s]; + for(; j+c<=sizeof(x->qs); j+=c) + for(size_t n=0;n<5;++n) for(size_t m=0;mqs[j+m]*pow3[n]; out[o++]=(float)((int)(((uint16_t)q*3)>>8)-1)*x->d; } + } + for(size_t n=0;n<4;++n) for(size_t h=0;h<2;++h){ + uint8_t q=x->qh[h]*pow3[n]; out[o++]=(float)((int)(((uint16_t)q*3)>>8)-1)*x->d; } +} + +int main(void) { + unsigned seed=99; + auto rnd=[&](){ seed=seed*1103515245u+12345u; return (seed>>16)&0xFFFF; }; + double worst_rel = 0.0; long checks = 0; long int_bad = 0, int_checks = 0; double worst_scaled = 0.0; + + for (int trial=0; trial<5000; ++trial) { + block_ptq1_0 w; + for (int i=0;i<24;++i) w.qs[i]=rnd()&0xFF; + for (int i=0;i<2;++i) w.qh[i]=rnd()&0xFF; + w.d = 0.01f + (rnd()%1000)/50000.0f; + + block_q8_1 y[4]; + for (int b=0;b<4;++b) { + y[b].ds.lo = 0.005f + (rnd()%1000)/80000.0f; y[b].ds.hi = 0.f; + for (int i=0;i<32;++i) y[b].qs[i]=(int8_t)((int)(rnd()%255)-127); + } + + float wf[QK_PTQ1_0]; ref_dequant(&w, wf); + + // EXACT test: the kernel's integer accumulator per chunk must equal the + // reference integer sum of trit*q8. This isolates logic from float rounding. + for (int c = 0; c < 4; ++c) { + int ref_sumi = 0; + for (int i = 0; i < 32; ++i) { + const int trit = (int) llround((double) wf[c*32+i] / (double) w.d); + ref_sumi += trit * (int) y[c].qs[i]; + } + // recompute the kernel's sumi by dividing its float result back out + const float got = vec_dot_ptq1_0_q8_1(&w, y, 0, c); + const int got_sumi = (int) llround((double) got / ((double) w.d * (double) y[c].ds.lo)); + if (ref_sumi != got_sumi) { ++int_bad; if (int_bad < 4) + printf(" INT MISMATCH trial %d chunk %d: ref %d got %d\n", trial, c, ref_sumi, got_sumi); } + ++int_checks; + } + + // reference dot over all four chunks, in float + double ref_total = 0.0; + for (int c=0;c<4;++c) + for (int i=0;i<32;++i) + ref_total += (double)wf[c*32+i] * ((double)y[c].qs[i] * (double)y[c].ds.lo); + + // kernel dot, chunk by chunk as MMVQ calls it + double got_total = 0.0; + for (int c=0;c<4;++c) got_total += vec_dot_ptq1_0_q8_1(&w, y, 0, c); + + // scale by the sum of magnitudes so cancellation in ref_total cannot inflate it + double mag = 0.0; + for (int c=0;c<4;++c) for (int i=0;i<32;++i) + mag += fabs((double)wf[c*32+i] * (double)y[c].qs[i] * (double)y[c].ds.lo); + const double denom = fabs(ref_total) > 1e-9 ? fabs(ref_total) : 1.0; + const double rel = fabs(got_total - ref_total) / denom; + if (rel > worst_rel) worst_rel = rel; + if (mag > 0) { const double sc = fabs(got_total - ref_total)/mag; if (sc > worst_scaled) worst_scaled = sc; } + ++checks; + } + printf(" dot products compared : %ld (4 chunks each)\n", checks); + printf(" worst relative error : %.3e\n", worst_rel); + printf(" worst err / sum|terms| : %.3e (immune to cancellation)\n", worst_scaled); + printf(" exact integer checks : %ld, mismatches %ld\n", int_checks, int_bad); + if (int_bad == 0) printf(" LOGIC EXACT: integer accumulator matches reference on every chunk\n"); + else printf(" LOGIC BUG in the kernel\n"); + return int_bad != 0; +} diff --git a/tests/test-ptq1_0-element-map.cpp b/tests/test-ptq1_0-element-map.cpp new file mode 100644 index 000000000000..d8cf77e10549 --- /dev/null +++ b/tests/test-ptq1_0-element-map.cpp @@ -0,0 +1,64 @@ +// Verifies the CUDA device accessor's element mapping matches the CPU codec exactly. +// The CUDA kernels cannot be compiled here (no nvcc), but this is where a mapping bug +// would hide, and it is pure integer logic, so it is checkable on the host. +#include +#include +#include +#include + +#define QK_PTQ1_0 128 +struct block_ptq1_0 { uint8_t qs[24]; uint8_t qh[2]; uint16_t d; }; + +// --- transcribed verbatim from ggml-cuda/dequantize.cuh --- +static int ptq1_0_trit(const block_ptq1_0 * x, const int e) { + uint8_t b; int n; + if (e < 80) { b = x->qs[e & 15]; n = e >> 4; } + else if (e < 120) { const int t = e - 80; b = x->qs[16 + (t & 7)]; n = t >> 3; } + else { const int t = e - 120; b = x->qh[t & 1]; n = t >> 1; } + uint32_t v = b; + for (int i = 0; i < 4; ++i) if (i < n) v = (v * 3) & 0xFF; + return (int)((v * 3) >> 8) - 1; +} + +// --- the CPU reference traversal from ggml-quants.c dequantize_row_ptq1_0 --- +static void cpu_ref(const block_ptq1_0 * x, int * out) { + const uint8_t pow3[6] = {1,3,9,27,81,243}; + const size_t stages[3] = {32,16,8}; + int o = 0; size_t j = 0; + for (size_t s = 0; s < 3; ++s) { + const size_t c = stages[s]; + for (; j + c <= sizeof(x->qs); j += c) + for (size_t n = 0; n < 5; ++n) + for (size_t m = 0; m < c; ++m) { + uint8_t q = x->qs[j+m] * pow3[n]; + out[o++] = (int)(((uint16_t)q * 3) >> 8) - 1; + } + } + for (size_t n = 0; n < 4; ++n) + for (size_t h = 0; h < sizeof(x->qh); ++h) { + uint8_t q = x->qh[h] * pow3[n]; + out[o++] = (int)(((uint16_t)q * 3) >> 8) - 1; + } +} + +int main(void) { + unsigned seed = 12345; + long bad = 0, total = 0; + for (int trial = 0; trial < 20000; ++trial) { + block_ptq1_0 blk; + for (int i = 0; i < 24; ++i) { seed = seed*1103515245u+12345u; blk.qs[i] = (seed>>16)&0xFF; } + for (int i = 0; i < 2; ++i) { seed = seed*1103515245u+12345u; blk.qh[i] = (seed>>16)&0xFF; } + int ref[QK_PTQ1_0]; cpu_ref(&blk, ref); + for (int e = 0; e < QK_PTQ1_0; ++e) { + ++total; + if (ptq1_0_trit(&blk, e) != ref[e]) { + if (bad < 5) printf(" MISMATCH blk%d e=%d cuda=%d cpu=%d\n", trial, e, ptq1_0_trit(&blk,e), ref[e]); + ++bad; + } + } + } + printf(" checked %ld element positions across 20000 random blocks\n", total); + printf(" mismatches: %ld\n", bad); + printf(" %s\n", bad==0 ? "CUDA element mapping MATCHES the CPU codec exactly" : "MAPPING BUG"); + return bad != 0; +} diff --git a/tests/test-quantize-fns.cpp b/tests/test-quantize-fns.cpp index 2811507a3cd2..4f77fb9701c7 100644 --- a/tests/test-quantize-fns.cpp +++ b/tests/test-quantize-fns.cpp @@ -160,6 +160,7 @@ static int test_vec_dot_q(bool verbose) { type == GGML_TYPE_TQ2_0 ? MAX_QUANTIZATION_TOTAL_ERROR_TERNARY : type == GGML_TYPE_Q2_0 ? MAX_QUANTIZATION_TOTAL_ERROR_TERNARY : type == GGML_TYPE_PQ2_0 ? MAX_QUANTIZATION_TOTAL_ERROR_TERNARY : + type == GGML_TYPE_PTQ1_0 ? MAX_QUANTIZATION_TOTAL_ERROR_TERNARY : type == GGML_TYPE_Q2_K ? MAX_QUANTIZATION_TOTAL_ERROR_2BITS : type == GGML_TYPE_IQ2_S ? MAX_QUANTIZATION_TOTAL_ERROR_2BITS : type == GGML_TYPE_Q3_K ? MAX_QUANTIZATION_TOTAL_ERROR_3BITS : @@ -185,7 +186,7 @@ static int test_vec_dot_q(bool verbose) { ? MAX_DOT_PRODUCT_ERROR_LOWBIT : type == GGML_TYPE_Q1_0 ? MAX_DOT_PRODUCT_ERROR_BINARY - : type == GGML_TYPE_TQ1_0 || type == GGML_TYPE_TQ2_0 || type == GGML_TYPE_Q2_0 || type == GGML_TYPE_PQ2_0 + : type == GGML_TYPE_TQ1_0 || type == GGML_TYPE_TQ2_0 || type == GGML_TYPE_Q2_0 || type == GGML_TYPE_PQ2_0 || type == GGML_TYPE_PTQ1_0 ? MAX_DOT_PRODUCT_ERROR_TERNARY : type == GGML_TYPE_NVFP4 ? MAX_DOT_PRODUCT_ERROR_FP4 diff --git a/tools/quantize/quantize.cpp b/tools/quantize/quantize.cpp index ff54010f247a..fe668c3ee123 100644 --- a/tools/quantize/quantize.cpp +++ b/tools/quantize/quantize.cpp @@ -35,6 +35,7 @@ static const std::vector QUANT_OPTIONS = { { "Q1_0", LLAMA_FTYPE_MOSTLY_Q1_0, " 1.125 bpw quantization", }, { "Q2_0", LLAMA_FTYPE_MOSTLY_Q2_0, " 2.25 bpw quantization (group 64)", }, { "PQ2_0", LLAMA_FTYPE_MOSTLY_PQ2_0, " 2.13 bpw quantization (group 128, Prism)", }, + { "PTQ1_0", LLAMA_FTYPE_MOSTLY_PTQ1_0, " 1.75 bpw ternarization (group 128, Prism)", }, { "Q4_0", LLAMA_FTYPE_MOSTLY_Q4_0, " 4.34G, +0.4685 ppl @ Llama-3-8B", }, { "Q4_1", LLAMA_FTYPE_MOSTLY_Q4_1, " 4.78G, +0.4511 ppl @ Llama-3-8B", }, { "MXFP4_MOE",LLAMA_FTYPE_MOSTLY_MXFP4_MOE," MXFP4 MoE", },