rocm: vectorized HIP PTQ1_0 vec_dot by Ankk98 · Pull Request #211 · PrismML-Eng/llama.cpp

Ankk98 · 2026-09-19T09:56:52Z

Overview

Through this branch I am trying to optimize Bonsai 2 27B to run efficiently on strix-halo through rocm path.

The model file stores weights as densely packed base-3 digits (trits). To use them, the GPU must unpack each trit and multiply-accumulate it with the activation. The old AMD code did this one trit at a time in a long serial chain (128 slow steps per block). The new code unpacks 4 trits at once using the GPU's byte-shuffle and 4-wide dot-product instructions. Same math, ~3x less instruction work, which is why decode got ~88% faster.

A wrong guess about a shuffle instruction's argument order was caught by on-device probes and unit tests before it could do harm; the corrected convention is now documented in a code comment.

Additional information

Tested L2 13593/13593
Deterministic end to end outout verified as compared to parent branch
Tested cross-arch compile (gfx90a/1030/1100/1151)

Bench summary (AC + high-perf, `-ngl 99 -fa on -ub 2048`, cooled between runs)

test	before (scalar)	after (vec)
tg128	11.47	21.0 - 21.5 (+~88%)
tg32	9.29	20.3 - 21.3 (+~120%)
tg128 ub512	8.89	19.90 (+124%)
pp512	175.4	186 - 204 (+6-16%)

Test Hardware

OS: Fedora Linux 44 (Workstation Edition) x86_64
Host: ROG Flow Z13 GZ302EA_GZ302EA (1.0)
Kernel: Linux 7.2.5-200.fc44.x86_64
CPU: AMD RYZEN AI MAX 390 (24) @ 5.06 GHz
GPU: AMD Radeon 8050S Graphics [Integrated]

Notes

Taken assistance of Opencode & Muse 1.3

Requirements

I have read and agree with the contributing guidelines
AI usage disclosure: YES

The model file stores weights as densely packed base-3 digits (trits). To use them, the GPU must unpack each trit and multiply-accumulate it with the activation. The old AMD code did this one trit at a time in a long serial chain (128 slow steps per block). The new code unpacks 4 trits at once using the GPU's byte-shuffle and 4-wide dot-product instructions. Same math, ~3x less instruction work, which is why decode got ~88% faster. A wrong guess about a shuffle instruction's argument order was caught by on-device probes and unit tests before it could do harm; the corrected convention is now documented in a code comment.

bri-prism

Agent review: posted by the maintainer's coding agent at their request.

No findings in this source pass. The four-byte trit extraction, element mapping and integer activation-sum correction are consistent with the previous scalar HIP path.

HIP compilation and device numerical tests were not run here. Before merging, retain a comparison against the scalar path that covers all packed-byte values and the qh tail on the supported AMD targets.

Reviewed commit: e435472d5391720bda720fbbbebf8d75938f8fca.

Copilot

Copilot was unable to review this pull request because the user who requested the review has reached their quota limit.

Copilot

Copilot review overview

🟡 Changes recommended

The architecture-specific vector path lacks regression coverage for its byte-permute and lane-packing behavior.

Get a fresh assessment by requesting another Copilot review.

Review effort: Balanced
Findings: 1 · 1

Open (2)

Add tests for vectorized trit packing and selector ordering · New
Keep comment sentences on one line and avoid “NB” · New

Ankk98 · 2026-09-21T09:03:35Z

+            const int q = (int) __builtin_amdgcn_perm(w_hi, w_lo, 0x07050301);
+            sumi[e >> 5] += ggml_cuda_dp4a(q, u, 0);
+            sumu[e >> 5] += ggml_cuda_dp4a(0x01010101, u, 0);


Addressed the comment, extended tests/test-ptq1_0-cuda-dot.cpp with:

a host transcription of the HIP vector idiom, unsigned {0,1,2} digits, bias folded into activation sums, qh tail in the unsigned scheme;

a hip_perm emulator encoding the device-proven rule, including the non-obvious part I verified while writing it: selectors are byte-granular ((s >> 8*i) & 7), not nibbles, nibble indexing contradicts the [b0,0,b1,0] comments and fails everything, byte indexing passes everything;

per-trial vec-vs-reference integer checks plus an exhaustive sweep (all 256 byte values × 6 packed words + 2 qh bytes × 4 chunks = 8192 exact checks).

Results: 20000 exact integer checks, 0 mismatches; 8192 exhaustive, 0 mismatches; exit 0. Device side stays covered by test-backend-ops on gfx1151.

Ankk98 · 2026-09-21T09:03:59Z

+    // Four packed bytes advance in the low bytes of 16-bit lanes, so one
+    // 32-bit multiply steps four trit streams at once (3*255 < 2^16, no
+    // cross-lane carry). Digits come out as {0,1,2}; the -1 folds into the
+    // activation sums, which a per-byte subtract cannot do borrow-free.
+    // NB: __builtin_amdgcn_perm picks sel bit2=0 from the SECOND arg, the
+    // opposite of CUDA __byte_perm, so the operands are swapped vs _multi.


Transcribe the vector idiom (unsigned digits, bias fold, swapped-operand perm, qh tail) and assert it against the scalar reference over all 256 byte values at every word position. Unwrap the vecdotq comment to one sentence per line. Assisted-by: Opencode with Muse Spark

bri-prism

Re-reviewed at 85a6c5f. The vectorization is correct on my read, the test is a real gate, and there is one evidentiary gap plus a sequencing hazard with #215.

The arithmetic checks out. I traced it by hand rather than trusting the commit message. Folding the bias as sum(q*a) - sum(a) with unsigned {0,1,2} digits is algebraically equivalent to the signed form. The element mapping is preserved: the scalar path's e = t*16 + m with m = 4w + j gives the same index the vector path computes as e = t*16 + 4w + j. The 0x07050301 selector places digit(b0..b3) in result bytes 0..3, which matches the little-endian order get_int_b4 loads the activations in. All three loops together touch each of the 128 elements exactly once, 80 plus 40 plus 8. No overflow is possible: sumu peaks near 32127 and sumi near 232*127, both far inside int32. The (e & 31) >> 2 group index stays 4-element aligned in both main loops. No findings there.

The byte-granular selector discovery is the valuable part of this PR and I am glad it is written down. Selectors being (s >> 8*i) & 7 rather than nibble-indexed, with sel bit2 = 0 choosing the second operand, is the kind of thing that costs someone a day.

Where I think the evidence is thinner than it looks.

hip_perm in the test is a host emulation of __builtin_amdgcn_perm, and vec_dot_ptq1_0_q8_1_vec is a transcription of the kernel rather than the kernel. That means the suite establishes that the algorithm is self-consistent given a belief about the intrinsic. It cannot establish the belief.

This matters more than usual here because of how the emulation was arrived at. You describe trying nibble indexing, finding it fails everything, then adopting byte indexing because it passes everything. That is fitting the emulator until the algorithm agrees with the reference. It is a reasonable way to find the rule, but it means the emulator and the kernel now share whatever assumption makes them agree, and the host test can no longer discriminate between "this matches the hardware" and "these two copies match each other". If the real intrinsic differs from hip_perm in some case neither of us has thought of, every one of those 28192 checks still passes and the device still returns wrong numbers.

The one artifact that would settle it is the device run, and that is the one reported without numbers. The host results are given precisely, 20000 exact integer checks and 8192 exhaustive, 0 mismatches, exit 0. The device side is given as "stays covered by test-backend-ops on gfx1151", which is an assertion rather than a result. Could you post the actual test-backend-ops output for the PTQ1_0 cases at 85a6c5f, with the case count? My earlier note asked for validation on the supported AMD targets, and gfx1151 output would close it outright.

Two smaller things.

The exhaustive sweep sets all four bytes of a packed word to the same value via v * 0x01010101u, so a within-word byte transposition such as b0 against b1 would pass that loop unnoticed. The 5000 random trials do use independent bytes and compare against ref_dequant, so the case is covered overall and I am not asking for a change. It is only that the exhaustive loop alone proves less than its name suggests.

The test carries its own copy of the kernel body. A later edit to the HIP branch in vecdotq.cuh will not be caught unless someone remembers to edit the transcription too. A short comment at the top of the HIP branch pointing at tests/test-ptq1_0-cuda-dot.cpp would make that link visible to whoever touches it next.

Sequencing with #215. Both PRs modify vec_dot_ptq1_0_q8_1, on opposite branches of the same #if. This PR leaves the #else calling vec_dot_ptq1_0_q8_1_multi, which is correct. #215 replaces that same #else with return 0.0f, which I have flagged there as a regression that silently zeroes PTQ1_0 on MUSA. Neither branch contains the other, and the two versions of this one function differ by about fifty lines, so they will conflict. Whoever resolves that conflict could reinstate the MUSA stub without noticing it is a defect, since it looks like an intentional part of the other change. Worth agreeing an order before either lands.

Method note: source reading and hand-tracing at 85a6c5f. I did not build, and I have no AMD hardware, so nothing above is a device result.

Ankk98 · 2026-09-22T20:24:04Z

Device code is byte-identical between e435472 and 85a6c5f (the delta is comment + host test only). So, this run covers both.

Full suite 13593/13593 green on the same device code
Host side 20000 exact integer checks, 8192 exhaustive, exit 0
Pointer comment linking the HIP branch to the test added in 3a82de9
On the transposition blind spot: acknowledged, no change — the 5000 random trials use independent bytes per position and check against ref_dequant, so that case is covered overall, as you note.

test-backend-ops test -b ROCm0 -j 16 -p 'type_a=ptq1_0' → 120/120 pass (skipped not supported combos excluded from the count):

ggml_cuda_init: found 1 ROCm devices (Total VRAM: 24576 MiB):
  Device 0: AMD Radeon 8050S Graphics, gfx1151 (0x1151), VMM: no, Wave Size: 32, VRAM: 24576 MiB
Testing 2 devices

Backend 1/2: ROCm0
  Device description: AMD Radeon 8050S Graphics
  Device memory: 24576 MB (24575 MB free)

  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=8,k=256,bs=[2,3],nr=[1,1],per=[0,1,3,2],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=1,b=1,m=512,n=5,k=256): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=1,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=1,k=128,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=8,n_used=4,b=0,m=512,n=1,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=8,n_used=4,b=1,m=512,n=5,k=256): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=1,k=256,bs=[3,2],nr=[2,2],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=2,b=1,m=512,n=5,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=4,b=1,m=512,n=4,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=8,n_used=2,b=0,m=512,n=17,k=256): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=2,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=1,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=4,k=256,bs=[3,2],nr=[2,2],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=1,b=1,m=512,n=17,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=8,n_used=4,b=1,m=512,n=17,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=4,b=1,m=512,n=5,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=8,n_used=1,b=0,m=512,n=129,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=2,b=1,m=512,n=17,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=8,n_used=4,b=0,m=512,n=4,k=256): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=8,k=256,bs=[2,3],nr=[1,1],per=[0,3,2,1],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=8,n_used=2,b=0,m=512,n=32,k=256): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=1,n=64,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=16,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=3,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=1,b=1,m=512,n=32,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=2,b=1,m=512,n=32,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=8,n_used=1,b=1,m=512,n=1,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=8,n_used=4,b=1,m=512,n=32,k=256): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=16,k=256,bs=[2,3],nr=[1,1],per=[0,2,1,3],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=8,n_used=4,b=0,m=512,n=5,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=4,b=1,m=512,n=17,k=256): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=16,k=256,bs=[1,1],nr=[2,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=8,n_used=2,b=0,m=512,n=129,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=2,b=0,m=64,n=16,k=384): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=4,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=1,b=1,m=512,n=129,k=256): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=16,k=256,bs=[1,1],nr=[1,2],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=1,b=0,m=512,n=1,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=8,n_used=1,b=1,m=512,n=4,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=8,n_used=2,b=1,m=512,n=1,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=8,n_used=4,b=0,m=512,n=17,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=2,b=1,m=512,n=129,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=4,b=1,m=512,n=32,k=256): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=5,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): ggml_backend_cuda_graph_compute: CUDA graph warmup complete
ggml_backend_cuda_graph_compute: CUDA graph warmup complete
ggml_backend_cuda_graph_compute: CUDA graph warmup complete
ggml_backend_cuda_graph_compute: CUDA graph warmup complete
OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=16,k=256,bs=[2,3],nr=[1,1],per=[0,1,3,2],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=8,n_used=4,b=1,m=512,n=129,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=2,b=0,m=512,n=1,k=256): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=16,k=256,bs=[3,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=6,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=8,n_used=2,b=1,m=512,n=4,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=1,b=0,m=512,n=4,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=8,n_used=1,b=1,m=512,n=5,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=4,b=0,m=512,n=1,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=8,n_used=4,b=0,m=512,n=32,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=2,b=0,m=512,n=1,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=4,b=1,m=512,n=129,k=256): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=16,k=256,bs=[2,3],nr=[1,1],per=[0,3,2,1],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=2,b=0,m=512,n=4,k=256): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=16,k=256,bs=[3,1],nr=[2,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=7,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=8,n_used=2,b=1,m=512,n=5,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=1,b=0,m=512,n=5,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=4,b=0,m=512,n=4,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=8,n_used=1,b=1,m=512,n=17,k=256): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=1,k=1024,bs=[3,2],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=2,b=0,m=512,n=32,k=256): OK
  OUT_PROD(type_a=ptq1_0,type_b=f32,m=256,n=1,k=1,bs=[1,1],nr=[1,1],trans_b=0): not supported [ROCm0] 
  OUT_PROD(type_a=ptq1_0,type_b=f32,m=256,n=1,k=1,bs=[1,1],nr=[1,2],trans_b=0): not supported [ROCm0, CPU] 
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=8,n_used=4,b=0,m=512,n=129,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=8,n_used=1,b=0,m=512,n=1,k=256): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=16,k=256,bs=[3,2],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=8,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=2,b=0,m=512,n=5,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=8,n_used=2,b=1,m=512,n=17,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=1,b=0,m=512,n=17,k=256): ggml_backend_cuda_graph_compute: CUDA graph warmup reset
ggml_backend_cuda_graph_compute: CUDA graph warmup reset
ggml_backend_cuda_graph_compute: CUDA graph warmup reset
OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=8,k=1024,bs=[3,2],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=8,n_used=1,b=1,m=512,n=32,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=4,b=0,m=512,n=5,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=8,n_used=4,b=1,m=512,n=1,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=8,n_used=1,b=0,m=512,n=4,k=256): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=16,k=256,bs=[3,2],nr=[2,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=9,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=2,b=0,m=512,n=17,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=8,n_used=2,b=1,m=512,n=32,k=256): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=16,k=1024,bs=[3,2],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=1,b=0,m=512,n=32,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=4,b=0,m=512,n=17,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=8,n_used=4,b=1,m=512,n=4,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=8,n_used=1,b=1,m=512,n=129,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=8,n_used=1,b=0,m=512,n=5,k=256): ggml_backend_cuda_graph_compute: CUDA graph warmup complete
ggml_backend_cuda_graph_compute: CUDA graph warmup complete
ggml_backend_cuda_graph_compute: CUDA graph warmup complete
ggml_backend_cuda_graph_compute: CUDA graph warmup complete
OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=16,k=256,bs=[3,2],nr=[1,2],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=1,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=2,b=0,m=512,n=32,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=8,n_used=2,b=1,m=512,n=129,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=4,b=0,m=512,n=32,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=1,b=0,m=512,n=129,k=256): ggml_backend_cuda_graph_compute: CUDA graph warmup reset
ggml_backend_cuda_graph_compute: CUDA graph warmup reset
ggml_backend_cuda_graph_compute: CUDA graph warmup reset
ggml_backend_cuda_graph_compute: CUDA graph warmup reset
OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=16,k=256,bs=[3,2],nr=[2,2],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=8,n_used=2,b=0,m=512,n=1,k=256): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=1,k=256,bs=[1,1],nr=[2,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=8,n_used=1,b=0,m=512,n=17,k=256): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=8,k=256,bs=[1536,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT(type_a=ptq1_0,type_b=f16,m=16,n=1,k=256,bs=[1,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): not supported [ROCm0, CPU] 
  MUL_MAT(type_a=ptq1_0,type_b=f16,m=16,n=1,k=256,bs=[1,1],nr=[2,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): not supported [ROCm0, CPU] 
  MUL_MAT(type_a=ptq1_0,type_b=f16,m=16,n=1,k=256,bs=[1,1],nr=[1,2],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): not supported [ROCm0, CPU] 
  MUL_MAT(type_a=ptq1_0,type_b=f16,m=16,n=1,k=256,bs=[3,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): not supported [ROCm0, CPU] 
  MUL_MAT(type_a=ptq1_0,type_b=f16,m=16,n=1,k=256,bs=[3,1],nr=[2,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): not supported [ROCm0, CPU] 
  MUL_MAT(type_a=ptq1_0,type_b=f16,m=16,n=1,k=256,bs=[3,2],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): not supported [ROCm0, CPU] 
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=2,b=0,m=512,n=129,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=1,b=1,m=512,n=1,k=256): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=1,k=256,bs=[1,1],nr=[1,2],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=1,k=256,bs=[2,3],nr=[1,1],per=[0,2,1,3],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=8,n_used=2,b=0,m=512,n=4,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=4,b=0,m=512,n=129,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=8,n_used=1,b=0,m=512,n=32,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=1,b=1,m=512,n=4,k=256): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=2,b=1,m=512,n=1,k=256): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=1,k=256,bs=[3,1],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=4,b=1,m=512,n=1,k=256): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=1,k=256,bs=[2,3],nr=[1,1],per=[0,1,3,2],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=8,n_used=2,b=0,m=512,n=5,k=256): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=1,k=256,bs=[3,1],nr=[2,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT_ID(type_a=ptq1_0,type_b=f32,n_mats=4,n_used=2,b=1,m=512,n=4,k=256): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=1,k=256,bs=[2,3],nr=[1,1],per=[0,3,2,1],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=1,k=256,bs=[3,2],nr=[1,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=1,k=256,bs=[3,2],nr=[2,1],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=4,k=256,bs=[2,3],nr=[1,1],per=[0,3,2,1],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=1,k=256,bs=[3,2],nr=[1,2],per=[0,1,2,3],k_v=0,o=1,src_overlap=0): OK
  MUL_MAT(type_a=ptq1_0,type_b=f32,m=16,n=8,k=256,bs=[2,3],nr=[1,1],per=[0,2,1,3],k_v=0,o=1,src_overlap=0): OK
  120/120 tests passed
  Backend ROCm0: OK
Backend 2/2: CPU
  Skipping
2/2 backends passed
OK

Sequencing with #215. Both PRs modify vec_dot_ptq1_0_q8_1, on opposite branches of the same #if. This PR leaves the #else calling vec_dot_ptq1_0_q8_1_multi, which is correct. #215 replaces that same #else with return 0.0f, which I have flagged there as a regression that silently zeroes PTQ1_0 on MUSA. Neither branch contains the other, and the two versions of this one function differ by about fifty lines, so they will conflict. Whoever resolves that conflict could reinstate the MUSA stub without noticing it is a defect, since it looks like an intentional part of the other change. Worth agreeing an order before either lands.

For the conflicts among the branches, we can merge the #211 first (it's hand-verified and device-proven, #215 is still open against prism), then #215 can rebase onto it. We need to inform the branch owner and make sure that the conflict resolver must preserve the vec_dot_ptq1_0_q8_1_multi call in the #else, the return 0.0f variant must not survive the merge. We can loop in the #215 owner i.e. @professorpalmer

@bri-prism Is this ok?

bri-prism

Agent review: posted by the maintainer's coding agent at their request.

Approving at 3a82de9. Thanks for posting the device output. The 120/120 PTQ1_0 test-backend-ops pass on gfx1151 checks __builtin_amdgcn_perm against the CPU reference on real hardware, which is the evidence I was asking for. The only change since 85a6c5f is the pointer comment, so that run covers this head.

On sequencing: yes, merging this first and then rebasing #215 onto it works for me, with one condition for whoever resolves the conflict. #215 (and #221, which carries it) now widens this branch to #if defined(GGML_USE_HIP) || defined(GGML_USE_MUSA). This PR puts __builtin_amdgcn_perm inside that branch, and it is an AMD-only builtin, so keeping #215's condition over this body would break the MUSA build. After the rebase, the vectorized path needs to stay under GGML_USE_HIP only, with MUSA on the scalar loop, and the #else must not end up returning 0.0f for any backend that can reach it. @professorpalmer, flagging this for the #215 rebase.

Method note: source reading at 3a82de9 plus your posted device output. I did not build or run it.

professorpalmer · 2026-09-23T03:57:42Z

Understood, and thank you for spelling out the condition.

After #211 is on prism, I will rebase #215 onto it with this split in vec_dot_ptq1_0_q8_1:

GGML_USE_HIP only: this PR's vectorized body. __builtin_amdgcn_perm does not compile on MUSA, so it cannot stay inside #if defined(GGML_USE_HIP) || defined(GGML_USE_MUSA).
GGML_USE_MUSA: the scalar loop from cuda: PTQ1_0 decode on consumer Ampere/Ada: SoA q8 activations, warp-per-row small-K GEMV (+10.5% TG, +14.7% with #220) #215.
#else: vec_dot_ptq1_0_q8_1_multi. The return 0.0f stub does not survive.

#221 carries the same function, so the same split applies there on its next rebase. I will not rebase either branch onto this one before it lands.

Implement PTQ1_0 conversion, row lookup, mat-vec and MMQ dispatch. Decode packed trits directly into MMQ tiles, use the full Ampere tile table and vectorize HIP dot products. Retain the MMQ batch cap and explicit Q8 activation precision. Source: PrismML-Eng#148 Source: PrismML-Eng#160 Source: PrismML-Eng#164 Source: PrismML-Eng#214 Source: PrismML-Eng#211 Source commits (PrismML-Eng/llama.cpp): 01fd952 9294043 7f292be bdc23b5 3b19c37 Co-authored-by: Cary Palmer <24235924+professorpalmer@users.noreply.github.com> Co-authored-by: Ankit Khandelwal <ankk98@gmail.com> Assisted-by: OpenAI Codex Assisted-by: GitHub Copilot Assisted-by: Opencode with Muse Spark Ported and rebased by Codex.

…ing, and -fa off (ggml-org#24503) * Add interface is_model_splitted() to check the c-graph is splited or not * Infer and propagate dynamic-dimension indices for all tensors in the GGML graph in api compute_model_outputs() * Only do this for fallback sub graph * Move dynamic dims compute in graph missmatch * ggml-openvino: fix tensor data handling for PERMUTE/VIEW ops in split models * ggml-openvino:add comments * ggml-openvino: override VIEW op_case to 0 for split model inputs * openvino backend: Handle unsupported VIEW shape-mismatch in OpenVINO backend * Enable additional mul_mat tests and add tensor data saving function (PrismML-Eng#81) * ggml-openvino: fix CONT/TRANSPOSE mapping and improve dynamic-dimension handling * OpenVINO: add NORM/TANH support and rework SOFT_MAX translation * ggml-openvino: extend VIEW handling * Enable -fa off (PrismML-Eng#118) * Enable --context-shift * Fix llm param compute error for normal softmax not the softmax in attention * OpenVINO backend: fix error for attention size compute in llm param * use tensor->extra in infer_request i/o * OpenVINO backend: refacter the compute_llm_params() func add get_attention_pattern_case to easy extand * OpenVINO backend: clean unused code * 1to1 match op update (PrismML-Eng#146) * added translate_1to1_match_1_input function and updated gelu and tanh translations * Remove unused translation function calls --------- Co-authored-by: Mustafa Cavus <mustafacavus@intel.com> * initial gemma4 support * removed hardcoded names for kv cache slicing * OpenVINO backend: Add new attention pattern for llm parameters compute * flash attn Q shape static conversion * Remove slice in permute translation when n_seq is 1 * return optional in extract_layer_from_name * OpenVINO backend: refactor VIEW related operation (PrismML-Eng#148) * OpenVINO backend: refactor VIEW related operation * Enable VIEW handling in following ops * OpenVINO backend does not support GGML_OP_NORM & GGML_OP_L2_NORM with VIEW input accuracy issue from OpenVINO * OpenVINO backend: Add ops l2_norm & pad * OpenVINO backend does not support CPY with non-contiguous data or mismatched types * add op SSM_CONV GATED_DELTA_NET * OpenVINO backend: fix error for bf16 in OV gpu plugin * reverted static Q input shape for attention layer * OpenVINO backend: remove hardcode name inp_tokens, which ignore some leaf case * Disable remote tensor due to bug in ov gpu * Disable n_token > 1 GATED_DELTA_NET on gpu * OpenVINO backend: fix the view op dynamic handling issue in gemma4 & enable view + get_row * OpenVINO backend: clean code * OpenVINO backend: enable view + norm/rms_norm * OpenVINO backend: concat op * OpenVINO backend: argsort op * OpenVINO backend: enable unary + view & GGML_UNARY_OP_SOFTPLUS * Fix issue for test-backend-ops in TOPK_MOE, which compare VIEW ops result, VIEW node in OpenVINO no need compare, the whole graph result is correct * OpenVINO backend: enable sum_rows * OpenVINO backend: enable clamp * OpenVINO backend: enable DIV * OpenVINO backend: enable GGML_OP_MUL_MAT_ID * OpenVINO backend: disable MUL_MAT_ID_FUSION case with large mem needed * OpenVINO backend: Disable GGML_OP_ARGSORT, cause test_backend-ops failed * OpenVINO backend: fix issue in mul_mat_id * OpenVINO backend: Disable DIV with broadcast on GPU * OpenVINO backend: update DIV * use ov internal op GatedDeltaNet * OpenVINO backend: enable llama erch test qwen3next * OpenVINO backend: enable RMS_NORM + VIEW & remove op_case 2 for rope * OpenVINO backend: fix error * suggested changes, need review * suggested changes, need review * OpenVINO backend: clean unused code & fix build warning * OpenVINO backend: enable minicpm3 for arch test * Disable GDN op (PrismML-Eng#177) * disable gated_delta_net * update stateful_kv_size correctly in mismatch case * OpenVINO backend: enable arch test for qwen3vl * OpenVINO backend: enable cohere2 for arch test * OpenVINO backend: enable t5 for arch test * OpenVINO backend: enable jamba for arch test * OpenVINO backend: remove warning for tmp * OpenVINO backend: enable kimi-linear for arch test * Remove unused * Fix gpt-oss accuracy issue * OpenVINO backend: enable arctic for arch test * OpenVINO backend: enable grok for arch test * Gemma4 initial npu support (PrismML-Eng#179) * Initiall gemma4 npu support * temp. fix for gemma4 accuracy bug on npu * Remove hardcoded names for npu-fold handling * revert static n tokens for cont translation as it is not needed * removed unused variable * ggml-openvino: add GGML_OPENVINO_ENABLE_CACHE env var to control decoder cache. Add environment variable GGML_OPENVINO_ENABLE_CACHE (default: YES). When set to NO, the decoder_cache is bypassed and models are rebuilt from the cgraph on every inference call in both dynamic and static compute paths. This is useful for debugging and verifying correctness without caching interference. * Revert "Gemma4 initial npu support (PrismML-Eng#179)" This reverts commit 0d29a9c. * OpenVINO backend: disable debug log print * Update TBB discovery. Delegated to OpenVINOs own config. * OpenVINO backend: GGML_OPENVINO_ENABLE_CACHE YES -> 1 * OpenVINO backend: fallback FLASH_ATTN_EXT in gemma3n to CPU backend * Add raw ov infer profiling metric * Add OV raw infer time metric to static compute path Co-authored-by: virajwad <84867530+virajwad@users.noreply.github.com> * Modify precision of static profiling * update to OV 2026.2, add OV windows CI * fix editorconfig-checks * Initiall gemma4 npu support * temp. fix for gemma4 accuracy bug on npu * Remove hardcoded names for npu-fold handling * revert static n tokens for cont translation as it is not needed * removed unused variable * test-llama-archs fix * Fix gemma4 flash_attn fallback * support im2col * fix code style * disable add_rope_sin_cos optimization * stateless boradcast and rope optimizations * Enable manual gqa attn by default for stateless gpu * manual gqa: fixed static batch * gemma4 llama-bench ctx update fix * Update OV win CI * stateful rope fusion temp. fix * OpenVINO backend: Conslolidate supported ops * Exclude unsupported GGML_OP_SUB cases * Exclude unsupported TOPK_MOE cases * OpenVINO Backend: MUL_MAT enhancements * Update OV CI * support f16 mask input for npu * Make GGML_OPENVINO_* env vars usage uniform Standardize all GGML_OPENVINO_* env flags: positive integers >0 to enable. Unset, empty, =0, or non-numeric values to disable. This fixes cases where text values or empty strings enabled features. * OpenVINO backend: Enhance envvar handling * more cleanup * move ggml_openvino_env_flag to appropriate place * OpenVINO backend: add REPEAT translator, Q5_1 weights, and GLU view-input fix * ggml-openvino: fix -Werror=cast-qual in extract_q5_1_data * Update openvino.Dockerfile Use BuildKit cache mounts for faster Docker rebuilds. Use apt instead of dpkg, remove unused .ddeb downloads, add DLLAMA_BUILD_TESTS=OFF. * ggml-openvino: centralize env var access via *getenv_str/getenv_int helpers Replace getenv and legacy flags with _str and _int helpers.Minor cleanup, doc updates. * OpenVINO backend: Enable GGML_OP_ADD_ID * Uptade openvino backend clamg-format * clang-format * Update OPENVINO.md (PrismML-Eng#211) * OpenVINO backend: fix accuracy issue for op CONCAT with i64 precision * Remove strict concurrency for gpu-openvino-low-perf * Update openvino CI keynames; add ccache-clear * Apply suggestions from code review Co-authored-by: Sigbjørn Skjæret <1629204+CISC@users.noreply.github.com> * Fix formatting --------- Co-authored-by: Xuejun Zhai <Xuejun.Zhai@intel.com> Co-authored-by: Mustafa Cavus <mustafa.cavus@intel.com> Co-authored-by: Mustafa Cavus <mustafacavus@intel.com> Co-authored-by: Xuejun <XuejunZhai@intel.com> Co-authored-by: Wang Yang <yang4.wang@intel.com> Co-authored-by: Ravi Panchumarthy <ravi.panchumarthy@intel.com> Co-authored-by: virajwad <84867530+virajwad@users.noreply.github.com> Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com> Co-authored-by: Mostafa Faheem <mostafaaafaheem@gmail.com> Co-authored-by: Sigbjørn Skjæret <1629204+CISC@users.noreply.github.com>

Implement PTQ1_0 conversion, row lookup, mat-vec and MMQ dispatch. Decode packed trits directly into MMQ tiles, use the full Ampere tile table and vectorize HIP dot products. Retain the MMQ batch cap and explicit Q8 activation precision. Source: PrismML-Eng#148 Source: PrismML-Eng#160 Source: PrismML-Eng#164 Source: PrismML-Eng#214 Source: PrismML-Eng#211 Source commits (PrismML-Eng/llama.cpp): 01fd952 9294043 7f292be bdc23b5 3b19c37 Co-authored-by: Cary Palmer <24235924+professorpalmer@users.noreply.github.com> Co-authored-by: Ankit Khandelwal <ankk98@gmail.com> Assisted-by: OpenAI Codex Assisted-by: GitHub Copilot Assisted-by: Opencode with Muse Spark Assisted-by: Claude (Opus 5.5), rebase onto ggml-org master with conflict resolution Ported and rebased by Codex.

github-actions Bot added ggml CUDA labels Sep 19, 2026

Ankk98 changed the title ~~rocm: vectorized HIP PTQ1_0 vec_dot [WIP]~~ rocm: vectorized HIP PTQ1_0 vec_dot Sep 19, 2026

bri-prism reviewed Sep 19, 2026

View reviewed changes

khosravipasha requested a balanced review from Copilot September 19, 2026 19:58

Copilot started reviewing on behalf of khosravipasha September 19, 2026 19:58

View session

Copilot AI reviewed Sep 19, 2026

khosravipasha requested a balanced review from Copilot September 21, 2026 00:32

Copilot started reviewing on behalf of khosravipasha September 21, 2026 00:32

View session

Copilot AI reviewed Sep 21, 2026

View reviewed changes

github-actions Bot added the testing label Sep 21, 2026

bri-prism reviewed Sep 21, 2026

View reviewed changes

ggml-cuda : point PTQ1_0 HIP branch at its host test mirror

3a82de9

bri-prism mentioned this pull request Sep 22, 2026

cuda: PTQ1_0 decode on consumer Ampere/Ada: SoA q8 activations, warp-per-row small-K GEMV (+10.5% TG, +14.7% with #220) #215

Closed

bri-prism approved these changes Sep 23, 2026

View reviewed changes

bri-prism merged commit 3b19c37 into PrismML-Eng:prism Sep 23, 2026
3 checks passed

professorpalmer mentioned this pull request Sep 25, 2026

cuda: Bonsai 2 27B at the full 262k window with the MTP head on 12 GB Ada: #218 + #215 hybrid PTQ1_0 dispatch, in-place q4_0/q8_0 K/V flash attention (integration checkpoint) #221

Merged

Navigation Menu

Repository navigation

rocm: vectorized HIP PTQ1_0 vec_dot - #211

rocm: vectorized HIP PTQ1_0 vec_dot#211
bri-prism merged 3 commits into
PrismML-Eng:prismfrom
Ankk98:strix-halo-optimization

Ankk98 commented Sep 19, 2026 •

edited

Loading

Uh oh!

bri-prism left a comment

Uh oh!

Copilot AI left a comment

Uh oh!

Copilot AI left a comment

Uh oh!

Ankk98 Sep 21, 2026

Uh oh!

Ankk98 Sep 21, 2026

Uh oh!

bri-prism left a comment

Uh oh!

Ankk98 commented Sep 22, 2026

Uh oh!

bri-prism left a comment

Uh oh!

professorpalmer commented Sep 23, 2026

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

Conversation

Ankk98 commented Sep 19, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Overview

Additional information

Bench summary (AC + high-perf, -ngl 99 -fa on -ub 2048, cooled between runs)

Test Hardware

Notes

Requirements

Uh oh!

bri-prism left a comment

Choose a reason for hiding this comment

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Copilot review overview

🟡 Changes recommended

Uh oh!

Ankk98 Sep 21, 2026

Choose a reason for hiding this comment

Uh oh!

Ankk98 Sep 21, 2026

Choose a reason for hiding this comment

Uh oh!

bri-prism left a comment

Choose a reason for hiding this comment

Uh oh!

Ankk98 commented Sep 22, 2026

Uh oh!

bri-prism left a comment

Choose a reason for hiding this comment

Uh oh!

professorpalmer commented Sep 23, 2026

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

Ankk98 commented Sep 19, 2026 •

edited

Loading

Bench summary (AC + high-perf, `-ngl 99 -fa on -ub 2048`, cooled between runs)