Skip to content

Read each K/V byte once in gqa-8 decode attention - #4077

Merged
zcbenz merged 3 commits into
ml-explore:mainfrom
dudududukim:du-sdpa-gqa
Aug 18, 2026
Merged

zcbenz merged 3 commits into
ml-explore:mainfrom
dudududukim:du-sdpa-gqa

Conversation

@dudududukim

@dudududukim dudududukim commented Aug 8, 2026 •

Copy link
Copy Markdown
Contributor

Closes #4076.

Proposed changes

sdpa_vector_2pass_1 has every simdgroup of a threadgroup stream the same K/V block, so loads scale with query heads instead of kv heads and gqa-8 models decode well below the machine's bandwidth (195 GB/s for 32q/4kv/hd128 and 120 GB/s for 64q/8kv/hd64 at 32K context, against 266 GB/s for a dense qmv on the same M5 Pro).

This adds sdpa_vector_2pass_1_gqa: each simdgroup owns a contiguous token sub-chunk of the block and computes HPT of its group's query heads from registers, so each K/V byte is read gqa_factor / HPT times instead of gqa_factor. The simdgroups' partials are merged once through threadgroup memory at the end; the partials layout and sdpa_vector_2pass_2 are unchanged. HPT is a template parameter because register pressure decides it: 8 heads per simdgroup wins at head_dim 64 but spills and loses at head_dim 128, where 4 is right.

Dispatch is deliberately narrow: single-token queries without mask or sinks, gqa_factor == 8, head_dim 64 or 128, N >= 8192. Everything else stays on the existing kernel. Sinks and other gqa factors can follow if this direction looks good.

All numbers below are from an M5 Pro 24 GB on AC, bf16, with K/V passed as strided cache slices and distinct buffers cycled between calls so the SLC doesn't flatter either arm. Both arms run in one process with the dispatch flipped between them (through a temporary env check that is not part of this diff), interleaved ABBA over several rounds; the tables give the median.

Kernel level:

geometry ctx main this PR speedup
32q/4kv/hd128 8192 90.8 us (185 GB/s) 84.8 us (198 GB/s) 1.07x
32q/4kv/hd128 16384 176.9 us (190 GB/s) 162.4 us (207 GB/s) 1.09x
32q/4kv/hd128 32768 343.5 us (195 GB/s) 301.5 us (223 GB/s) 1.14x
64q/8kv/hd64 8192 147.1 us (114 GB/s) 127.7 us (131 GB/s) 1.15x
64q/8kv/hd64 16384 288.5 us (116 GB/s) 247.5 us (136 GB/s) 1.17x
64q/8kv/hd64 32768 561.6 us (120 GB/s) 440.6 us (152 GB/s) 1.27x

Per-round spread is within 0.01x of the median for every row above.

End-to-end decode, ms/step, on Qwen3-30B-A3B-4bit truncated to 24 layers (so 32K of KV fits in 24 GB; 32q/4kv/hd128) and on the full 36-layer Qwen2.5-3B-Instruct-4bit (16q/2kv/hd128). Context 4096 is below the dispatch gate, so both arms run the identical kernel there and it serves as a control:

ctx 30B main 30B this PR delta 3B main 3B this PR delta
4096 (control) 7.289 7.292 -0.0% 9.626 9.626 -0.0%
8192 8.596 8.301 +3.5% 10.451 10.178 +2.7%
16384 10.757 9.936 +8.3% 12.152 11.479 +5.9%
32760 15.025 13.726 +9.5% 15.674 14.355 +9.2%

Individual control rounds span -0.2%..+0.3% (30B) and -0.6%..+1.6% (3B), so the gated contexts sit well outside the noise floor.

Logits stay within bf16 rounding of the current kernel with argmax stable over greedy decode, and outputs match an fp32 reference to 1e-5 (fp16) and 2e-4 (bf16) on both gated and ungated shapes. The added test covers both instantiated head dims at a context that is not a multiple of the block count.

Checklist

  • I have read the CONTRIBUTING document
  • I have run pre-commit run --all-files to format my code / installed pre-commit prior to committing changes
  • I have added tests that prove my fix is effective or that my feature works
  • I have updated the necessary documentation (if needed)

@zcbenz zcbenz added the await verification This pull request is non-trivial and requires a human expert to verify its correctness. label Aug 9, 2026
@erwinzhang7

Copy link
Copy Markdown
Contributor

Verified on an M5 Max, 40 GPU cores, 128 GB, macOS 26.2. Your numbers are from an
M5 Pro 24 GB, so this is a different part rather than a bigger one, and the 128 GB
removes the truncation your 30B run needed.

base throughout is 6a0dd0fd2, this branch's own parent, rather than current
main, since the branch carries a merge commit and diffing against main would
credit it with three weeks of unrelated work.

The kernel win holds on wider silicon. bf16, K/V as
strided cache slices with eight buffer sets cycled so the SLC cannot hold them,
12 rounds of 20 calls with one eval per round:

geometry ctx base this PR speedup yours
32q/4kv/hd128 8192 53.4 us 46.2 us 1.16x 1.07x
32q/4kv/hd128 16384 87.6 us 80.9 us 1.08x 1.09x
32q/4kv/hd128 32768 156.2 us 139.3 us 1.12x 1.14x
64q/8kv/hd64 8192 73.5 us 61.4 us 1.20x 1.15x
64q/8kv/hd64 16384 130.2 us 109.4 us 1.19x 1.17x
64q/8kv/hd64 32768 244.0 us 186.6 us 1.31x 1.27x

The baseline on this machine is 1.7x to 2.3x faster than yours on the same kernel
before any change. The worry with a bandwidth argument is that it stops applying on a part with
bandwidth to spare. The argument holds: larger on four of the six
geometries, and within a couple of points on the other two.

Nothing outside the gate moves at all. Both builds scored against each other
and against an fp32 reference:

case max abs diff vs base vs fp32
ungated, ctx 4096 below the gate 0 3.5e-4
ungated, gqa_factor 4 0 3.3e-4
ungated, head_dim 80 0 4.6e-4
ungated, 2 query tokens 0 3.4e-4
gated 32q/4kv/hd128 2.4e-4 2.8e-4
gated 64q/8kv/hd64 2.4e-4 1.9e-4
gated, ctx 8192+137 4.9e-4 2.4e-4

Bit-identical on every ungated shape, so the dispatch condition fires exactly
where you say and nowhere else. On the gated shapes the differences are bf16
rounding and the new kernel sits as close to fp32 as the old one, so it isn't
sacrificing accuracy for speed. The 8192+137 row would expose an off-by-one in the sub-chunk
partitioning, and it is clean.

End to end, on all 48 layers rather than 24. Same model, Qwen3-30B-A3B-4bit,
untruncated, 128 GB holds 32K of KV for the whole thing. 24 decode steps per
context after a full prefill, median:

ctx base this PR delta yours, 24 layers
4096 9.303 ms 9.260 ms +0.5% -0.0%
8192 10.682 ms 10.264 ms +3.9% +3.5%
16384 12.757 ms 12.177 ms +4.5% +8.3%
32760 17.724 ms 16.491 ms +7.0% +9.5%

The win is real at full depth and clears the control comfortably. My 4096 control
comes out at +0.5% rather than zero (the two arms are separate processes with
separate model loads) so read 0.5% as this harness's noise floor.

However, the magnitude is smaller than the truncated run shows, and consistently:
+4.5% against +8.3% at 16K, +7.0% against +9.5% at 32K. That's from halving the
layer count: it removes half the non-attention work per step, so attention
occupies a larger share of what is left and any speedup measured against the step
scales up with it.

I'd quote +7.0% at 32760 on the full model rather than +9.5%. It is still a
good number for a decode path and it is one nobody can argue with. Good work!

@zcbenz zcbenz left a comment •

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Can you add the affected shapes to sdpa_bench.py?

Sorry just realized this is sdpa_vector, ignore this comment.

@zcbenz zcbenz removed the await verification This pull request is non-trivial and requires a human expert to verify its correctness. label Aug 18, 2026
if (!mask && !sinks && q.shape(2) == 1 && gqa_factor == 8 &&
q.shape(-1) == v.shape(-1) && (q.shape(-1) == 64 || q.shape(-1) == 128) &&
N >= 8192) {
std::string gname;

@zcbenz zcbenz Aug 18, 2026 •

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It seems that most of the host code is duplicated, I think you can just change the kname?

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Right!
Fixed in d1ab9de, the gate now just appends _gqa to the kname.
Tests pass and the dispatch still switches at the boundary.

@zcbenz zcbenz left a comment

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Nice improvement, thanks!

@zcbenz
zcbenz merged commit fa0d446 into ml-explore:main Aug 18, 2026
28 checks passed
davidtai added a commit to Layr-Labs/mlx that referenced this pull request Aug 25, 2026
* Return tuple in meshgrid (ml-explore#4229)

* Add endpoint parameter to linspace (ml-explore#4184)

Co-authored-by: Cheng <git@zcbenz.com>

* Fix vmap of partition/argpartition dropping the kth argument (ml-explore#4116)

* Fix nan_to_num replacing inf with 0 for float16 and bfloat16 (ml-explore#4222)

Co-authored-by: codeAnqiang-ma <273298913+codeAnqiang-ma@users.noreply.github.com>
Co-authored-by: Cheng <git@zcbenz.com>

* Fix einsum not broadcasting batch dimensions in batched tensordot (ml-explore#4125)

Co-authored-by: Cheng <git@zcbenz.com>

* Dequantize in float32 (ml-explore#4241)

* chore: Reject complex in erf and erfinv (ml-explore#4243)

* Fix cpu compilation failure of abs with uint (ml-explore#4240)

Co-authored-by: Cheng <git@zcbenz.com>

* Fix quantize matrix multiplication floor issue (ml-explore#4251)

* Only use MPI backend for world size > 1 (ml-explore#4210)

* chore: Reject complex in expm1, sigmoid and arctan2 (ml-explore#4257)

* Decompose small kernel-depth 3D convs into 2D convs (ml-explore#3785)

Co-authored-by: katlun-lgtm <264247399+katlun-lgtm@users.noreply.github.com>
Co-authored-by: Cheng <git@zcbenz.com>

* Fix Metal sort of a view with a negative stride (ml-explore#4252)

* Mirror the depth axis in the decomposed 3D conv when flipped (ml-explore#4277)

* Fix Metal row reductions on negative-stride views (ml-explore#4267)

Co-authored-by: Fu Xiaonan <214359569+FU-max-boop@users.noreply.github.com>

* [CUDA] Fix custom kernel cache collision for same name, different source (ml-explore#4273)

Co-authored-by: Cheng <git@zcbenz.com>

* Fix ops rejecting integers larger than INT32_MAX (ml-explore#4255)

Co-authored-by: Feli <feli@hnu.edu.cn>
Co-authored-by: Cheng <git@zcbenz.com>

* Fix var/std for complex numbers (ml-explore#4260)

* Fix int32 overflow in conv padded input and pad shapes (ml-explore#4258)

Co-authored-by: Cheng <git@zcbenz.com>

* chore: Reject complex in remainder (ml-explore#4270)

* chore: Compare the macOS SDK version as a version when gating JACCL (ml-explore#4286)

* Clamp ring socket transfers so a payload of 2 GiB or more can be sent (ml-explore#4281)

Co-authored-by: Cheng <git@zcbenz.com>

* chore: Use normalize_axis_index in split/unstack/partition/topk (ml-explore#4288)

* Remove grouped output in CI (ml-explore#4195)

* [CUDA] Fix finding cuda 13 headers in JIT compilation (ml-explore#3995)

* Refactor wheel building script (ml-explore#3818)

* Make mx.compile cache erasing thread safe (ml-explore#4248)

Co-authored-by: yentur <mr.yentur@gmail.com>

* Add builds for free-threaded python (ml-explore#3812)

* Fix int32 overflow in concatenate/repeat/kron (ml-explore#4303)

* python: Widen list elements that do not fit in int32 to int64 (ml-explore#4305)

* Propagate CPU errors to events (ml-explore#3742)

Co-authored-by: Alessio Pollero <alessio.pollero@gmail.com>

* Fix mx.arange dtype inference overflow regression (ml-explore#4324)

* Add workflow to update pull request limit bypass list (ml-explore#4320)

* Support head dimension 72 in Metal full attention (ml-explore#4330)

* Patch bump to 0.32.2 (ml-explore#4333)

* Preserve subnormal float values when casting to bool (ml-explore#4224)

* python: Support assigning through a bare Ellipsis index (ml-explore#4314)

* Fix divmod truncating the quotient for floats (ml-explore#4108)

Co-authored-by: Cheng <git@zcbenz.com>

* Add force_fused option to scaled_dot_product_attention (ml-explore#4185)

* chore: Reject negative eps in the normalization layers (ml-explore#4312)

* Bound GGUF metadata string/array values against the file mapping (ml-explore#4212)

Co-authored-by: x14ngch3n <x14ngch3n@users.noreply.github.com>
Co-authored-by: Cheng <git@zcbenz.com>

* Read each K/V byte once in gqa-8 decode attention (ml-explore#4077)

* Fix fft vmap and jvp for transforms over a subset of axes (ml-explore#4138)

* Fix median dropping NaN (ml-explore#4146)

* Fix the CPU scan over a size one axis with a padded stride (ml-explore#4139)

Co-authored-by: Cheng <git@zcbenz.com>

* chore: Validate the optimizer betas at construction (ml-explore#4310)

Co-authored-by: Cheng <git@zcbenz.com>

* `RMSNormVJP` backward writes a full `{n_rows, D}` `gw_temp` intermediate (ml-explore#4293)

* [Bug]: add default none value to axis parameter of the take_along_axis (ml-explore#4357)

Co-authored-by: Anastasiia Filippova <a_filippova@apple.com>

* Add a fused full-attention path for head_dim 256 on NAX devices (ml-explore#3842)

Co-authored-by: Cheng <git@zcbenz.com>

* Update nanobind to 2.15.0 (ml-explore#4337)

* Skip unnecessary simdgroup computations for quantised MOE matmuls on NAX (ml-explore#4352)

* Add AI usage policy (ml-explore#4331)

Co-authored-by: Jake Bowhay <60778417+j-bowhay@users.noreply.github.com>

* Raise cpu stream errors from synchronize (ml-explore#4338)

Co-authored-by: Cheng <git@zcbenz.com>

* chore: Validate eps in Adam at construction (ml-explore#4361)

Co-authored-by: Anastasiia Filippova <a_filippova@apple.com>

* Bound winograd conv2d working set by tiling the batch (ml-explore#4102)

Co-authored-by: Cheng <git@zcbenz.com>

* Use a 32-row block in qmm_t_nax when one block covers all of M (ml-explore#4171)

* chore: Deduplicate fftshift and ifftshift (ml-explore#4318)

* Fix Log and Equal is_equivalent ignoring primitive state (ml-explore#4266)

Co-authored-by: Cheng <git@zcbenz.com>

* Stabilize reduced-precision InstanceNorm (ml-explore#4230)

* chore: Normalize negative axes in sort and argsort (ml-explore#4332)

* Clean up main thread compile cache before python interpreter shuts down (ml-explore#4373)

* chore: Check malformed jaccl hostfile that miss rdma in pairs (ml-explore#4284)

Co-authored-by: Cheng <git@zcbenz.com>

* Round mxfp8 block scales up to avoid saturation (ml-explore#4353)

Co-authored-by: Daniel Hiltgen <daniel.hiltgen@ollama.com>
Co-authored-by: Cheng <git@zcbenz.com>

* Add support for the __array_namespace_info__  (ml-explore#4334)

* Stop a failed CUDA graph commit from poisoning the encoder (ml-explore#4356)

Co-authored-by: Cheng <git@zcbenz.com>

* Fix quantized kernels in JIT build (ml-explore#4372)

Co-authored-by: Cheng <git@zcbenz.com>

* Avoid zero work in stride-2 ConvTranspose3d (ml-explore#4343)

* [CUDA] Ce fused kernel (ml-explore#3947)

* Fix cpu exclusive scan for complex numbers (ml-explore#4272)

Co-authored-by: Cheng <git@zcbenz.com>

* Support Relocatable CUDA DLLs on Windows (ml-explore#4382)

* Use cast_to for fused AsType in compiled Metal kernels (ml-explore#4351)

Co-authored-by: katlun-lgtm <katlun@windyviews.com>
Co-authored-by: Cheng <zcbenz@gmail.com>

* python: Declare DLPackCompatible protocol members as methods (ml-explore#4384)

* Fix quantizing sliced arrays (ml-explore#4381)

* Fix einsum dropping a trailing empty subscript (ml-explore#4299)

Co-authored-by: Cheng <git@zcbenz.com>

* Add script to run python tests (ml-explore#4393)

* Hold GIL in AttachedData destructor (ml-explore#4391)

* Bound Metal buffer COUNT, not just bytes, in MetalAllocator

The Metal allocator throws `[metal::malloc] Resource limit (N) exceeded`
when num_resources_ (the live+cached Metal buffer COUNT) reaches
resource_limit_ (the iogpu.rsrc_limit sysctl, default ~499000). Freed
buffers are recycled into a size-keyed cache whose only trim is by BYTES
(release_cached_buffers takes a bytes-to-free target, max_pool_size_ ~=
physical RAM). Under churn with many distinct buffer shapes (varied prompt
lengths, growing KV caches, multiple co-resident models) the cache fills
with entries never reused at that exact size, so the COUNT climbs to the
limit while byte usage stays modest and the byte trim never fires — the
process crashes mid-inference on a machine with most of its RAM free.

malloc() now also reclaims by count: when num_resources_ crosses a 90%
high-water mark of resource_limit_, it clears the (pure-reuse) buffer
cache so the count drops back to the live working set. Clearing the cache
only costs re-allocation, never correctness, so the count limit becomes
unreachable by any request mix or batching method while the existing byte
limits keep total memory bounded.

Adds get_num_resources()/get_resource_limit() to the public memory API
(metal + no_gpu + cuda backends) so the count and its ceiling are
observable from callers. Adds an MLX_RESOURCE_LIMIT env override that can
only LOWER the ceiling (clamped to the OS limit, strictly validated) to
exercise the trim deterministically and as an operator safety valve.

* perf(mlx): opt-in Gemma 4 expert-QMM tile kernel with parallel descriptor builder (#4)

* perf(mlx): add opt-in Gemma 4 expert-QMM tile kernel with parallel descriptor builder

Adds a distinctly-named expert QMM implementation for the Gemma 4
26B-A4B MoE production shapes, gated by MLX_GATHER_QMM_EXPERT_SLICES:

- qmm_t_expert_impl: BM32 expert tile body (BM16 fallback rows) taking a
  private/by-value row count; the shared qmm_t_impl constant-address ABI
  and all ordinary gathered/batched/dense QMM routes are unchanged.
- build_gemma4_sorted_expert_tiles_bm32: one 128-thread threadgroup
  replaces the reference design's single-GPU-thread serial builder;
  parallel expert-range binary search, Hillis-Steele scan, and strided
  upper-bound descriptor emission.
- Selector runs after the NAX-first route and requires affine BF16
  transposed inputs, 4-bit gs=64 weights, 128 experts, assignment counts
  of exactly 4096/8192/16384, and the exact gate/up or down rank-3
  shapes; every miss keeps the legacy route. NAX engagement is
  non-engagement, never bypassed.
- device.{h,cpp}: one-shot request resolution, nonthrowing dual-symbol
  AOT probe/prewarm, relaxed-atomic diagnostics (requested, aotAvailable,
  naxAvailable, hits, per-class fallbacks).
- gpu_tests: exact-shape arithmetic parity, fallback, and counter
  invariant probes.

Retention standing (2026-08-09 production matrix): opt-in experiment.
Standalone profile dropped (prefill -10.2% vs bracket); paired
weighted-unsort+R1 profile retained-final (prefill +1.8%, TTFT -7.5%,
decode +3.3%, arrival E2E +12.0%). NOTE: this source post-dates the
benchmarked binaries/metallib (post-measurement kernel-body edit);
rebuild and re-verify before any performance claim.

* fix(mlx): fail-safe sortedness check in gemma expert tile builder; counter/atomic hygiene

Review-wave fixes for the R1 expert-QMM path:

- N1 (sortedness trust): build_gemma4_sorted_expert_tiles_bm32 now
  verifies each thread's post-binary-search segment boundary against the
  generalized invariant indices[start - 1] < lid <= indices[start]
  (edge threads check their single neighbor), votes per simdgroup via
  simd_or, folds the votes through threadgroup memory, and on any
  violation retracts count[0] to 0 (tile kernel then early-returns) and
  records the violation in count[1]; the buffer ABI is unchanged
  (count index 1 was previously unused). try_gemma4_expert_qmm allocates
  the second count element, drains the encoder after the builder, and
  re-routes a retracted call to the order-agnostic legacy path instead of
  dispatching the tile kernel (zero count is unambiguous: the selector's
  assignment gate guarantees M is 4096/8192/16384).
- N2 (route-condition duplication): the sorted-RHS gate literal that
  appeared (negated) in the diagnostics record and in the dispatch
  decision is now the shared static constexpr predicate
  takes_sorted_rhs_route, so future tuning of the 16/4 thresholds cannot
  desynchronize counter vs route.
- N3 (per-call bias normalization): gather_qmm_rhs no longer spends
  ensure_row_contiguous on biases before classification reads the raw
  tensor's fields; normalization runs only inside the winning-route
  branch (hit semantics unchanged; the legacy block keeps its own
  normalization point and ordering).
- N4 (armed_ data race): Gemma4ExpertQMMCounters::armed_ is now
  std::atomic<bool> with relaxed loads/stores in armed(), snapshot(),
  snapshot_and_disarm() (read-then-write order preserved) and
  clear_and_arm(); the class remains non-copyable, now enforced.

* fix(mlx): make the R1 sortedness fail-safe sound; proper retract attribution

F1: the per-expert boundary vote was a partial detector -- an inversion
inside a segment used by no other expert's boundary could escape, so
"re-route on any violation" overclaimed. build_gemma4_sorted_expert_tiles_bm32
now also runs a strided adjacent-pair scan: thread lid checks
indices[i-1] <= indices[i] for i = lid+1; i < M; i += 128, covering every
adjacent pair in [1, M) exactly once (1..128 iterations at the reachable
M in {4096,8192,16384}). Adjacent-pair monotonicity is transitive, so a
clean scan is a sound and complete sortedness oracle; it folds into the
same simd_or/threadgroup vote and the same retract (count[0]=0, count[1]=1).
The boundary checks stay as cheap, precise diagnostics.

F2: retracts were write-only in count[1] and surfaced as
fallback_metallib_unavailable -- misattribution in the only observable
surface. A dedicated fallback_sortedness_retracted counter now rides the
GemmA4 route counters and the C diagnostics ABI
(sizeof 80 -> 88, new uint64 at offset 80; existing offsets unchanged).
try_gemma4_expert_qmm returns the route class: count[0]==0 with count[1]==1
records fallback_sortedness_retracted, any other unusable build keeps
fallback_metallib_unavailable, then re-routes to the legacy path as before.

F4: new doctest drives the full armed() -> clear_and_arm() ->
snapshot_and_disarm() cycle and the attempts == hits + fallbacks invariant
including the new class; the route-table and counter-invariant tests now
cover fallback_sortedness_retracted.

Verified: cmake tests 262/262 + 3550 assertions pass; metal -Wall -Wextra
-fno-fast-math compile of kernels/quantized.metal is warning-free.

* perf(metal): E=256 expert-tile route + trust + gpu::eval UAF fix — darkbloom-base mirror (#7)

* perf(metal): instantiate E=256 expert-tile route for Qwen 3.5/3.6 MoE prefill (mirror of Cmlx/mlx 58fab46)

* fix(metal): use-after-free in gpu::eval for primitives that synchronize mid-eval (mirror)

* perf(metal): trust mode skips retract readback (mirror)

* fix(compile): preserve all-cache binding cleanup

---------

Co-authored-by: JasonHonKL <148705846+JasonHonKL@users.noreply.github.com>
Co-authored-by: AK <144495202+AKnassa@users.noreply.github.com>
Co-authored-by: Cheng <git@zcbenz.com>
Co-authored-by: Adityaj0 <93090622+Adityaj0@users.noreply.github.com>
Co-authored-by: anchor <codeanqiang@gmail.com>
Co-authored-by: codeAnqiang-ma <273298913+codeAnqiang-ma@users.noreply.github.com>
Co-authored-by: Rohan Gautam <rohan1gautam@gmail.com>
Co-authored-by: Ayaan Gazali <ayaangazali.work@gmail.com>
Co-authored-by: Erwin Zhang <59893706+erwinzhang7@users.noreply.github.com>
Co-authored-by: katlun-lgtm <katlun@gmail.com>
Co-authored-by: katlun-lgtm <264247399+katlun-lgtm@users.noreply.github.com>
Co-authored-by: robertomeroni <150194833+robertomeroni@users.noreply.github.com>
Co-authored-by: Fu Xiaonan <ht3fudatou@163.com>
Co-authored-by: Fu Xiaonan <214359569+FU-max-boop@users.noreply.github.com>
Co-authored-by: Hao Xu <hxu44@apple.com>
Co-authored-by: Feli <89400571+FeliGame@users.noreply.github.com>
Co-authored-by: Feli <feli@hnu.edu.cn>
Co-authored-by: Eyüp Can Akman <eyupcanakman@gmail.com>
Co-authored-by: Cheng <zcbenz@gmail.com>
Co-authored-by: yentur <mr.yentur@gmail.com>
Co-authored-by: Alessio Pollero <alessio.pollero@gmail.com>
Co-authored-by: Zhiqi Zhang <zhiqizhangg@gmail.com>
Co-authored-by: Daniel Hiltgen <dhiltgen@users.noreply.github.com>
Co-authored-by: Tanish Jain <recklurker@gmail.com>
Co-authored-by: hojin12312 <hojin12312@gmail.com>
Co-authored-by: Xiang Chen <46052474+x14ngch3n@users.noreply.github.com>
Co-authored-by: x14ngch3n <x14ngch3n@users.noreply.github.com>
Co-authored-by: Duhyeon, Kim <49020301+dudududukim@users.noreply.github.com>
Co-authored-by: rohith <kapellirohith@gmail.com>
Co-authored-by: Ishaan Samantray <devteam.aegis@gmail.com>
Co-authored-by: Aaishwarya Mishra <aaishwarymishra@gmail.com>
Co-authored-by: Anastasiia Filippova <a_filippova@apple.com>
Co-authored-by: Yanzhao Wang <19340816+wyanzhao@users.noreply.github.com>
Co-authored-by: XXXXRT666 <157766680+XXXXRT666@users.noreply.github.com>
Co-authored-by: Jake Bowhay <60778417+j-bowhay@users.noreply.github.com>
Co-authored-by: vraj patel <87225460+vraj00222@users.noreply.github.com>
Co-authored-by: Gusanidas <33495733+Gusanidas@users.noreply.github.com>
Co-authored-by: Dwijen Patel <dwijen@gmail.com>
Co-authored-by: Vladimir Iglovikov <ternaus@users.noreply.github.com>
Co-authored-by: Brian C. <94733710+deBrian07@users.noreply.github.com>
Co-authored-by: Daniel Hiltgen <daniel.hiltgen@ollama.com>
Co-authored-by: YH Yan <strayberry0w0@gmail.com>
Co-authored-by: katlun-lgtm <katlun@windyviews.com>
Co-authored-by: anupsv <6407789+anupsv@users.noreply.github.com>
Co-authored-by: Gajesh Naik <26431906+Gajesh2007@users.noreply.github.com>
Co-authored-by: David Tai <davidtai@Davids-MBP.lan>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

sdpa_vector_2pass_1 loads K/V once per query head instead of once per kv head

3 participants