apps/ggml: GGML quantized kernels on Approximations by alexreinking · Pull Request #9511 · halide/Halide

alexreinking · 2026-10-06T17:47:33Z

Work in progress; top of stack #9654. Reimplements GGML's quantized vec_dot and mul_mat in Halide on Func::approximate_by() (#9486), and compares them against GGML's CPU paths (plain, repack, KleidiAI) on Neon. Metal comes later. Touches only apps/, except a temporary commit carrying #9655 and #9656 (dropped at the next restack). Replaces #9417.

Harness (ggml-quant-bench): GGML v0.26.0 through a vcpkg port, reached only through public headers. An f64 oracle with tolerances derived from each provider's declared precision; randomized provider order; medians with 95% CIs; --cold for weights streamed from DRAM.
Kernels: q4_0 and q8_0 schemes, and one codec generator (bit-exact against GGML). A single mul_mat algorithm, out(n, m) = sum_k W(k, n) * X(k, m), serves q8_0, f16 and f32 activations, with vec_dot, gemv and gemm as schedules. Weight layouts: GGML's plain records, its 4x4 and 4x8 repacked (AoSoA) records, and an SoA layout, all expressed as Approximations. On arm_i8mm targets the gemm uses smmla on 4x16, 4x8 or 4x4 register tiles whose accumulators split_storage keeps in registers; the gemv on 4x4 records uses by-element sdot (Test by-element sdot/udot codegen on ARM #9544). kernels/ is 477 lines, with no per-format code.

Results so far (M3 Max, AC power). vec_dot, 1 thread: GGML/Halide is 0.93 at n = 256 (adapter call overhead) and 0.99-1.00 at n = 4096 and 14336. mul_mat with q4_0 weights on 60 Llama 3, Llama 3.2 and Qwen 2.5 layer shapes, warm and cold, each side at its best thread count (GGML/Halide, >1 means Halide is faster; ranges, median in parentheses):

activations	GGML path	gemm (M = 8, 32, 512)	gemv (M = 1)
q8_0	plain (fastest)	1.52-2.45 (1.83)	0.98-1.86 (1.07)
f32	plain	1.38-2.03 (1.81)	0.95-1.55 (1.04)
f32	fastest (repack or KleidiAI)	0.64-1.02 (0.79)	0.64-1.25 (0.99)

Against plain GGML (the same weight layout), every row is faster or within the 95% CIs. The app links its own copy of Halide's thread pool (apps/ggml/runtime/), whose lock-free do_par_for fast path and keep-awake bring fork/join at 8-12 threads down from 22-40 us to 1-1.5 us.

With interleaved weights (preliminary; both AC runs were noisy, load 2-17): against GGML's fastest f32-activation path, all 24 gemv rows are at or above 1.0 (geomean 1.17); gemm geomean is 1.08 at M = 8 and 0.99 otherwise, with 7 rows at 0.93-0.97 (M = 8 and 32, against repack or KleidiAI).

Metal (on top of the CPU work): gemv rows for every block format (K-quants, i-quants, mxfp4), codebooks staged in threadgroup memory, and a simdgroup_matrix gemm for M >= 8 (f16) / M >= 16 (exact f32). The gemm uses #9283 (on main) and the Metal simdgroup_matrix backend of #9619 (in this stack). Weights run it on f16 operands as GGML's kernel_mul_mm (q4_0:f16, q8_0:f16, and f16 weights, which need #9655 and #9656).

Tests: ctest ggml_check (ggml-quant-bench --check), every GGML type against the oracle plus the Halide codecs and kernels, 0 failures; ggml_thread_pool, the vendored pool.

Authored by GitHub Copilot

macOS malloc page-aligns large blocks, so a K=4096 f32 act row (16 KB) ends exactly at a 16 KB page edge, and whether the next page was ever touched depends on the process's heap history. When it wasn't, the hardware prefetcher (running ~2 KB ahead of the act stream) walks the missing translation on every call -- failed walks aren't cached -- and long-bodied kernels that stream f32 acts directly (grid IQ, K-quants, tq, iq4_xs) ran 1.7-3.2x slower in a random subset of processes (IPC 5.5 -> 2.4, ~44 failed data-side table walks per call, 0 TLB fills). Allocate w and a with 16 KB of written slack past the end (skipped by --no-pretouch, alongside the image pretouch). Real activations live in reused compute buffers whose neighbouring pages are resident, so the fast mode is the representative one. 24 interleaved processes, vec_dot K=4096 f32, slow processes before -> after: iq2_xs 6 -> 0, iq2_xxs 9 -> 0, iq3_xxs 6 -> 0; mul_mat 4096x2048x1 T1 (20 processes): 7/10/8 -> 0/0/0. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

staged_indices() staged a table decode's index array in a GPR a GPR-wide word (8 bytes on arm64) at a time. q5's qh (4 bytes, a byte array in the algorithm since the precomputed luts) then loaded as an 8-byte word shifted inwards over the block's other fields: q5_0's from 2 bytes before the block (the previous block's qs, d, qh; for the first block, 2 bytes before the weights' buffer). Every index byte took a shift and a mask (not ubfiz), and LLVM took d from that word (lsr, fmov) instead of loading it. Stage an array narrower than a GPR as one word of its own (q5: ldr w at qh's offset, as GGML's memcpy to a uint32_t): the hot loops are the old u32 layout's again (q5_0 q8_0 vec_dot 67 -> 63 instructions per 2 blocks; vec_dot, gemv and gemm loops of q5_0/q5_1 the same as with the u32 layout modulo registers). Only q5_0/q5_1 vec_dot and mul_mat CPU objects change; Metal unchanged (bytes). M3 Max, 8 interleaved processes, vs the u32 layout (8fa2d0b): 0.98-1.01x on vec_dot (256, 4096; q8 and f32 acts) and mul_mat 4096x2048xM (M = 1, 4, 32; 1 and 10 threads); vs the bytes layout before: q5_0 q8_0 vec_dot 4096 0.86x, f32 0.96x, q5_1 q8_1 0.96x, gemv q5_0 0.93-0.99x. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

metal_gemm cloned dot into out (clone_in), specialized out for the simdgroup-matrix gemm, then rewrote that specialization's values to call dot again: a write to out's Definition from the schedule. Read dot through a wrapper instead (dot.in(out)), eager_inline it into the gemm's specialization alone, clone dot into the wrapper, then eager_inline the wrapper into out's other definitions: the gemm's branch reads dot, the rows read the clone, with public directives only. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Replace deferred_shifts(), the last schedule code that rewrote Func definitions, with schedule directives. hoisted_shifts() keeps the read-only analysis of where the 2-bit fields sit by chunk (runs and position classes). It then splits the chunks, rfactors the codes over the position class, bounds the lane and class, and calls hoist_invariants(), which now hoists the fields' invariant right shift as a factor: each field at bits 2 and 4 is read in place and its class's sum is shifted once; bit 6's (192) doesn't fit the int8 code and keeps its shift. stage_definition() is removed with it. Bit-exact: --check 2630/0. vec_dot objects are identical to before. The mul_mat (gemv path) objects differ only by the write-back stage's register allocation (a few movi/mov per loop), and the timing difference is within noise. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Move the measurement essays out of the schedule comments; keep the decision, the key number and the correctness constraints. Comment-only change. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

metal_rows' nr and sg and metal_flat's nr, sg and c were never overridden; make them locals. Tiles::tasks was an int used as a flag: make it bool balance. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

#9541 is in HEAD, so specializations no longer share bounds by RVar name; name the RVar ryi like the other recipes. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Build a ggml::Knobs from the GeneratorParams once in matmul.cpp instead of threading five to seven positional knobs through vec_dot, mul_mat, blocked and blocked_mmla. blocked derives its gemv-only knobs from a single gemv flag; row_bytes moves into Formats. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Replace eleven hand-rolled caller searches and name-based set differences with four small helpers. Codegen-neutral: every CPU and Metal object is byte-identical. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

The two recipes shared their prologue, the codes and quads chain, the per-block register storage, the act-codes staging and the accumulator epilogue. Two-level weights are now blocked_rows' sub-block case. Codegen-neutral: every CPU and Metal object is byte-identical. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Seven sites specialized a stage only when a condition was defined; they now call when(stage, cond). Codegen-neutral: every CPU and Metal object is byte-identical. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

inline_operands and inline_until were the same loop, differing in a filter (pure Funcs only) and whether f is in its own stop set. Both are now inline_callees(f, kept, pure). Codegen-neutral: every CPU and Metal object is byte-identical. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

vec_dot and mul_mat both vectorized a staged act of values after scheduling its encoder; encoder() now takes that vector width. Codegen-neutral: every CPU and Metal object is byte-identical. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

derived_sums now returns false (no fold) instead of throwing when hoist_invariants can't split a constant code offset off exactly, as with iq1_s's and iq1_m's delta. That drops their four fold_offsets=none entries from formats.cmake. Codegen-neutral: every object is identical apart from two local closure names in one iq1_m object. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Codegen-neutral: every CPU and Metal object matches the baseline (apart from the previous commit's two iq1_m closure names). Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

The linear layers of a 2B-parameter diffusion transformer at one denoising step's token counts (M = 1159 and 4231), and its AdaLN projection on ~20 rows. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

q8_0's Metal libraries ran their large-M gemm (metal_mm) on f32 operands, while GGML's kernel_mul_mm dequantizes q8_0 to half in threadgroup memory and multiplies half simdgroup fragments into f32 accumulators. Add the q8_0:f16 weight compute format row (f32:f16 and f16 acts), as q4_0:f16: each decoded weight rounded to f16 (Precision::w_quant), its library running metal_mm on f16 operands for M >= 8 and N >= 32. The harness's halide-by-m-q8_0_to_f16_* providers pick it for M > 16 and the exact q8_0 library below. Denoiser preset, ggml-metal time / best Halide time (M3 Max, 15 rounds; qkvo, up, down at M = 1159, the same at M = 4231, then adaln at M = 20): q8_0 f32: before .79 .78 .80 .79 .79 .77 .83 after .99 .97 .98 .97 .97 .96 .89 q8_0 f16: before .78 .76 .77 .78 .79 .77 .81 after .98 .96 .96 .96 .96 .96 .93 Small M (llama3-8b-full, M = 1, 8: by-m runs the unchanged exact library); at M = 32 by-m's f16 gemm is 1.04-1.13x faster than the exact library's f32 gemm (e.g. 4096 x 14336 f32 acts 0.513 vs 0.577 ms). --check --threads 4: 2648 results, 0 failures; ctest 10/10. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

metal_mm's threadgroup tile on f16 operands (the :f16 rows) grows from 64 x 32 to 64 x 64 at M >= 256, and from 16 x 16 to 32 x 16 at 16 < M < 32. The f32-operand gemm keeps its tiles: 64 x 64 at K steps of 64 would use 32 KB of threadgroup memory, and 32 x 16 measured 0.6-0.85x of 16 x 16 there. Denoiser preset, ggml-metal time / best Halide Metal time, M3 Max, 2 x 15 rounds interleaved, medians (qkvo, up, down at M = 1159, then at M = 4231, then adaln): q4_0 f16 before 0.93 0.92 0.92 0.93 0.92 0.91 0.84 after 0.95 0.97 0.95 1.00 0.99 1.00 0.87 q4_0 f32 before 0.95 0.95 0.96 0.95 0.94 0.93 0.83 after 0.94 0.95 0.92 0.96 1.00 1.00 0.89 q8_0 f16 before 0.98 0.96 0.98 0.96 0.94 0.96 0.95 after 0.97 0.98 0.96 1.03 1.01 1.01 0.98 q8_0 f32 before 0.99 0.95 0.97 0.97 0.96 0.98 0.87 after 0.95 0.98 0.97 1.02 1.00 1.01 0.96 At M = 512 (4096 x 4096, 4096 x 14336, 14336 x 4096) the :f16 rows run 1.02-1.12x faster than before. At 16 < M < 32, 32 x 16 is 1.0-1.14x at 2048 x 12288 x 20 and 0.97-1.02x at 4096 x 4096 x 20/24, 14336 x 4096 x 20. Tried and dropped (neutral or worse): ggml's 8 x 8-contiguous threadgroup layout, K steps of 32, 16/8-value staging chunks per lane (0.7x at adaln), m-fastest threadgroup order, and 128 x 32, 128 x 64, 64 x 128 and 128 x 64 (8 simdgroups) tiles at large M. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Squashed copies of #9655 (keep dense ramps dense when flattening aliased GPU allocations) and #9656 (settle comparisons using learned differences and their alignment), needed by the Metal gemm for f16 weights. Drop this commit when restacking onto a main that has them. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

DEPENDS ON CORE PRs #9655 (FlattenAliasedAllocations keeps Ramps dense) and #9656 (Simplify_LT settles comparisons with known_difference plus alignment). On the core without them the generator fails ("The tensor core fragment ... is used outside a tensor core instruction": the K steps' guards on the flat reduction, which #9656 proves away, keep a fragment load outside the MMA). #9655 makes its threadgroup stores vectors instead of one per lane. f16 weights ran metal_flat on Metal at every M: rows, no simdgroup_matrix, 0.36-0.50x of ggml-metal (which runs kernel_mul_mm on them for M > 8) at large M. Now f16 storage with f16 acts (f16, f32:f16) gives f16_operands, and metal_gemm runs metal_mm on them (K in steps of 64 or, if K % 64 != 0, 32 values; 64 x 32 tiles at large M, which beat 64 x 64 here: 0.93 vs 0.99 at M = 1159) where M >= 8, N >= 32, K % 32 == 0 and M * N >= 2^16, and metal_flat (with its own staged acts) otherwise, in the same kernel. Below M * N = 2^16 the rows win: 1.2-2.6x at M = 8..32, N = 512..2048. Denoiser preset, ggml-metal time / best Halide Metal time, M3 Max, core with #9655 + #9656, 2 x 15 rounds interleaved, medians (qkvo, up, down at M = 1159, then at M = 4231, then adaln): f16 f16 before 0.50 0.46 0.49 0.47 0.44 0.45 0.83 after 0.99 0.99 0.99 0.98 0.98 0.99 1.07 f16 f32 before 0.42 0.40 0.42 0.39 0.37 0.40 0.67 after 0.98 0.97 0.97 0.96 0.96 0.94 1.04 Shapes full (M = 1, 8, 32, 512; 12 shapes), f16 weights, new / old time: M = 1, 8 unchanged (geomean 0.99-1.01, every shape within 0.94-1.19x); M = 32 1.22x (f16 acts) and 1.36x (f32:f16) geomean, at worst 0.96x; M = 512 2.09x and 2.24x, at least 1.58x. The exact f32-act library (no f16 operands) is unchanged. bf16 keeps metal_flat (f16_operands takes Float(16) only). check 2648 / 0 fails, f16/bf16 Metal on shapes full 744 / 0, on the denoiser preset 56 / 0, ctest 10 / 10 (all against #9655 + #9656). Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

On f16 operands at 16 < M < 32, metal_mm ran two overlapping ShiftInwards 16-row tiles, so every weight block was decoded twice. Now one 32 x 32 tile covers the rows (as GGML's kernel_mul_mm does) and decodes the weights once. - The tile's dot is RoundUp on m. So the act staging computes rows past M and reads acts there. The spec now clamps those rows to the last act row: Xf for acts converted in the kernel (cn != an; previously rowwise only), and, after sever, a `_rows` wrapper of each act stand-in that every caller is redirected to. The schedule inlines it into the threadgroup staging, which gets one min() per half4 load, while the fragment loads stay plain. This is value-preserving: each output row depends only on its own act row. No core change and no new directive. - out reads the tile from a threadgroup copy (edge = dot.in(via), inlined in every other specialization), because a guarded store straight from Tile memory doesn't lower. Its m is split with PredicateStores, which leaves the calls unpredicated. A GuardWithIf there put edge's call under an if, and SkipStages then gave every gemm kernel uniform .used/.loaded branches in its K loop: f16 2048x2048x32 went from 49 to 66 us. The 32 x 16 f16 shape is dropped (unreachable now). Measured on M3 Max against ggml-metal, ratio = ggml / halide (by-m), denoiser preset, 2 x 25 rounds: q4_0 f16 acts 2048x12288 M=20: 0.91 -> 1.01 (209.5 -> 190.4 us) q4_0 f32 acts 2048x12288 M=20: 0.90 -> 0.99 (218.8 -> 201.5 us) q8_0 f16 / f32 acts: 0.98 -> 1.07 / 0.97 -> 1.03 M = 1159 and 4231: unchanged within noise (0.93-1.01). Small M (llama3.2-1b-full, M = 1, 8, 32, 512, every Metal provider): geo-mean time ratios 0.997-1.02. CPU kernels: 0.98-1.01, with the f32_to_q8_0 kernels (which now clamp in the act quantization) at 0.99-1.00. --check --threads 4: 2648 results, 0 failures. ctest: 10/10. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

… staging q4_0 "down" in the denoiser preset (K 7168, N 2048, M 1159) ran at 0.93x of ggml-metal while the other large-M shapes were at parity. An M sweep shows the per-row cost of the 64 x 64 tile is within 1-2% of GGML's kernel_mul_mm where there is no edge waste (M = 1024, 1088, 1280). At M = 1159 = 18 * 64 + 7 the 64 x 64 tiles compute 1216 rows (GGML's 64 x 32: 1184, +2.7%), and their 32 x 19 = 608 threadgroups put up to 16 on each of the M3 Max's 40 cores (95% balance; GGML's 1184 -> 30, 98.7%). A model of time ~ max threadgroups per core x tile height fits every tile variant tried within about 1%. metal_mm now picks a 64 x 56 tile (blocked f16 operands, M >= 256) where it computes fewer rows past M than 64 x 64 (M = 1159: 1176 rows). The rule only looks at edge waste, not the core count; simulated over M = 256..5000 for 10 to 40 cores it gains 0.5% on average. On f32 acts (f32:f16) the 56-row tile was 7-9% slower than the model, with generated code identical to f16's except for the loads. Unrolling the staging loops (each lane's 4-value pieces of ws and xs) fixes that, and gains about 1% on the 64 x 64 tile as well. The unroll applies to every metal_mm tile. The changes are schedule-only. There is no spec or core change. Ratio ggml/halide (by-m), 4efb479 -> this, denoiser preset, 2x15 rounds: q4_0 f16 down/1159 0.92 -> 1.03 (2x20-round repeat 0.93 -> 1.01) q4_0 f32 down/1159 0.92 -> 0.98 (repeat 0.93 -> 1.00) q4_0 qkvo/1159 0.93/0.94 -> 1.03/1.00, up/1159 0.97/0.96 -> 0.99/1.00 q8_0 1159 shapes 0.95-0.97 -> 1.00-1.05; 4231 shapes unchanged (0.99-1.02) f16 f16 (4x15 alternating rounds): 0.95-0.97 -> 0.99 on the 1159/4231 shapes; f16 f32:f16 geo 0.988 -> 0.989 (within noise) adaln unchanged (1.02-1.06) Small M (M = 1, 8, 32, 512, all Metal providers, every blocked format): within noise at M <= 32. At M = 512 the exact blocked libraries are 1-9% faster from the unrolled staging. CPU kernels are untouched. check 2648/0; app ctest 10/10. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…2.9) Only nvfp4 has float per-sub-block scales, and it sets no sublanes or gpr_scales knob and is not affine, so those three clauses never decide. Codegen-neutral: all 757 CPU and 110 Metal objects byte-identical. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

… 2.1) The 16-lane form (no split) was marked untimed. Timed A/B on M3, 15 alternating rounds: q6_K q8_K vec_dot 1.049-1.058x slower, gemv at one thread 1.055-1.063x (q8_K and f32:q8_K acts). Keep the split; comment only. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Retimed on M3 (A/B, 15 alternating rounds): without sublanes iq3_xxs's q8_K vec_dot is now faster (0.971x at K=14336, 0.975x at K=4096). Its gemv still gains from them (f32:q8_K at K=14336, one thread: 1.014x slower without, CI [1.011, 1.028]), so the two mul_mat entries stay. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

A/B on M3 (15 alternating rounds), the two q6_K row_pairs=false entries removed: gemv at K=4096 on one thread 1.020x (q8_K act, CI [1.018, 1.030]) and 1.029x (f32:q8_K, [1.014, 1.046]) slower; at K=14336 and on ten threads within noise (the old note's 0.85-0.95x on ten no longer holds). Keep the entries; comment only. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

The float chain's middle entry, 2 x 1 row pairs where N is 2 or 3 and M > 1, was slower than single rows. A/B on M3 (15 alternating rounds; f16, bf16 with f16/bf16/f32 acts; K=4096/14336, N=2/3, M=8..64; one and ten threads): 48 cases 0.41-0.81x the time (1.2-2.4x faster); the f16_acc16 tree libs, whose 2 x 2 tiles already cover N >= 2, within noise or slightly faster. The f16 gemv's interleave 8 (f79a80f) is unchanged. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Every one-level record with two records per gemv tile (the decoder-type check dropped) is within noise on short runs but 1.02-1.11x slower on q8_0.4x4/4x8 at K = 14336 on ten threads in sustained (200 ms) runs, so the check stays; its comment now records the retimed cost. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…m 2.3) Retimed every interleave=1 entry against the default 2 (Apple M3, 15 alternating rounds, T1/T10): q2_K's vec_dot is 1.05x faster with 2, q4_K's and q6_K's within noise or faster, iq1_s's gemv 1.02-1.05x faster. Their interleave=1 goes (two entries, two tokens). The ldur rule (interleave 1 on every q8_K act) is not adopted: it makes iq3_s's and iq1_m's vec_dot 1.14-1.18x slower and several gemvs 1.01-1.13x, and q3_K's vec_dot stays at 1 (1.02x slower with 2 in sustained runs). Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

alexreinking added this pull request to stack #9512 October 6, 2026 17:47

alexreinking changed the title ~~alexreinking/ggml quant~~ apps/ggml: GGML quantized kernels on Approximations Oct 6, 2026

alexreinking added the skip_buildbots Do not run buildbots on this PR. Must add before opening PR as we scan labels immediately. label Oct 6, 2026

alexreinking removed this pull request from stack #9512 October 6, 2026 19:33

alexreinking force-pushed the alexreinking/ggml-quant branch from 0542a85 to 3f8e310 Compare October 6, 2026 19:34

alexreinking added this pull request to stack #9522 October 6, 2026 19:34

alexreinking mentioned this pull request Oct 6, 2026

Add Func::approximate_by() and composable Approximations #9486

Draft

alexreinking force-pushed the alexreinking/ggml-quant branch 2 times, most recently from 1627d58 to 9ed2641 Compare October 6, 2026 20:16

alexreinking force-pushed the alexreinking/ggml-quant branch from 9ed2641 to 401c635 Compare October 7, 2026 06:39

alexreinking removed this pull request from stack #9522 October 7, 2026 16:08

alexreinking added this pull request to stack #9524 October 7, 2026 16:08

alexreinking force-pushed the alexreinking/ggml-quant branch from 097c0b7 to cb9cd78 Compare October 7, 2026 16:08

alexreinking force-pushed the alexreinking/ggml-quant branch from cb9cd78 to 68152e9 Compare October 7, 2026 17:26

alexreinking force-pushed the alexreinking/ggml-quant branch from 68152e9 to f631a55 Compare October 7, 2026 17:47

alexreinking removed this pull request from stack #9524 October 7, 2026 19:21

alexreinking force-pushed the alexreinking/ggml-quant branch from f631a55 to 74a1b3f Compare October 7, 2026 19:26

alexreinking added this pull request to stack #9527 October 7, 2026 19:26

alexreinking removed this pull request from stack #9527 October 8, 2026 11:09

alexreinking force-pushed the alexreinking/ggml-quant branch from 74a1b3f to 7bc3997 Compare October 8, 2026 11:10

alexreinking added this pull request to stack #9529 October 8, 2026 11:10

alexreinking and others added 4 commits November 2, 2026 18:48

alexreinking mentioned this pull request Nov 2, 2026

Add Stage::hoist_invariants(groups) and Stage::value(s) #9508

Open

4 of 5 tasks

alexreinking and others added 25 commits November 2, 2026 19:25

ggml: shorten schedule.h comments to decisions and key numbers

8c3bc57

Move the measurement essays out of the schedule comments; keep the decision, the key number and the correctness constraints. Comment-only change. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

ggml: make never-overridden schedule parameters constants

6a4bc9d

metal_rows' nr and sg and metal_flat's nr, sg and c were never overridden; make them locals. Tiles::tasks was an int used as a flag: make it bool balance. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

ggml: drop the ryi_rows workaround

674009b

#9541 is in HEAD, so specializations no longer share bounds by RVar name; name the RVar ryi like the other recipes. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

apps/ggml: add when() for an optional specialization

7fb391b

Seven sites specialized a stage only when a condition was defined; they now call when(stage, cond). Codegen-neutral: every CPU and Metal object is byte-identical. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

apps/ggml: use without() for the act decoders above the derived ports

8ff2806

Codegen-neutral: every CPU and Metal object matches the baseline (apart from the previous commit's two iq1_m closure names). Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

apps/ggml: add a denoiser shapes preset

c21a9e5

The linear layers of a 2B-parameter diffusion transformer at one denoising step's token counts (M = 1159 and 4231), and its AdaLN projection on ~20 rows. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Navigation Menu

Repository navigation

apps/ggml: GGML quantized kernels on Approximations - #9511

apps/ggml: GGML quantized kernels on Approximations#9511
alexreinking wants to merge 218 commits into
alexreinking/approximationsfrom
alexreinking/ggml-quant

alexreinking commented Oct 6, 2026 •

edited

Loading

Uh oh!

This branch has not been deployed

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

Conversation

alexreinking commented Oct 6, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

This branch has not been deployed

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

alexreinking commented Oct 6, 2026 •

edited

Loading