Skip to content

Add force_fused option to scaled_dot_product_attention - #4185

Merged
zcbenz merged 1 commit into
ml-explore:mainfrom
DT-HJ:force-fused-sdpa
Aug 18, 2026
Merged

zcbenz merged 1 commit into
ml-explore:mainfrom
DT-HJ:force-fused-sdpa

Conversation

@DT-HJ

@DT-HJ DT-HJ commented Aug 11, 2026 •

Copy link
Copy Markdown
Contributor

Proposed changes

Add a force_fused option to mx.fast.scaled_dot_product_attention that
ignores the automatic dispatch heuristics, uses a fused kernel, and raises
with the unsupported constraint when no fused kernel is available. This lets
inference runtimes — the only party that knows their resident weights, KV
cache size, and memory budget — decide when the fused path's bounded-memory
behavior is worth its throughput cost.

Also restores the fused full-attention support for head_dim=192/256 that was
proposed in #3293/#3660. Those kernels are now reachable only through
force_fused; the default dispatch does not route to them.

Follow-up hardening incorporates the full-dispatch threadgroup safety check
and capability-diagnostic concerns raised by @apocryphx in #4186, while
retaining this PR's 192/256 kernels and C++ positional-call compatibility.

Background

#3658 tracks models with genuine head_dim=192/256 full-attention layers. At
long context, the unfused path materializes a score
transient of O(n_heads × qL × kL) per full-attention layer. The fused
full-attention kernel tiles K/V so its transient stays bounded regardless of
kL. However, benchmark measurements on a genuine head_dim=256 workload
(the Qwen 3.6 35B checkpoint measured in #3658) show the fused path trades
prompt throughput for that memory floor:

context main prompt_tps fused prompt_tps main peak fused peak
32K 1005.743 871.489 41.316 39.882
64K 802.929 624.513 44.201 40.068

These stock mlx-lm benchmark arms were separated by other main-branch
commits, so they are supporting workload evidence rather than a strict
single-commit A/B. The optional 131K fused arm aborted with an
interactivity-impacting command-buffer error and is intentionally omitted.

A fixed dispatch threshold cannot pick correctly for all runtimes, so the
decision is exposed to the caller instead. This PR is the follow-up to the
maintainer request in #3658:

Let's add a force_fused=True option to mx.fast.scaled_dot_product_attention API that ignores the heuristics and uses fused path, and throws when a fused kernel is not available for the shape. — @zcbenz

Supersedes the closed #3660.
Addresses #3658.

API

mx.fast.scaled_dot_product_attention(
    q, k, v, *, scale, mask=None, sinks=None, force_fused=False, stream=None
)
  • force_fused=False (default): preserves existing dispatch for previously
    supported shapes. The missing head_dim=192 vector kernel is added but
    reachable only through force_fused.
  • force_fused=True: bypass the dispatch heuristics and use a fused kernel.
    Raises ValueError with the unsupported constraint when no fused kernel is
    available, which covers:
    • shapes not covered by the full- or vector-attention fused kernels
      (including unsupported head_dims, e.g. 512);
    • CPU execution;
    • training mode that requires a logsumexp output (the fused kernels do not
      emit one).

is_training on its own (when the VJP path already falls back) is allowed to
use the fused forward kernel — the output is identical.

Kernel support

  • Full-attention (steel_attention): add head_dim=192/256 instantiations.
    Selected only when force_fused=True. Extend the existing large-head-dim
    V-tile synchronization barriers from BD == 128 to BD >= 128, covering
    192/256 as well as 128. This synchronization gap was also independently
    identified in giaki3003/mlx@1175b4d.
  • Vector (decode) (sdpa_vector): add the missing head_dim=192 (192,192)
    and aggregation instantiations. The new 192 vector kernel is opt-in via
    force_fused like the new full-attention kernels; the existing head_dim=256
    vector kernel keeps its automatic routing.
  • Route head_dim >= 192 away from the NAX kernel family, which has no
    192/256 instantiations.
  • Check the compiled pipeline's maximum threadgroup size before both full
    attention dispatches, avoiding silent zero output on register-limited GPUs.
  • Use each backend's actual support matrix: Metal full/vector kernels and CUDA
    cuDNN/vector kernels share the same opt-in semantics.

Validation

  • Focused tests in python/tests/test_fast_sdpa.py:
    • force_fused=True full attention for head_dim 192/256 at short key
      lengths (where the default dispatch stays unfused) matches the reference;
    • force_fused=True full attention for head_dim 192/256 at kL=16385
      matches the reference, covering the Steel V-tile barrier path;
    • force_fused=True on an already-supported head_dim=128 matches;
    • force_fused=True on an unsupported shape (head_dim=512) raises;
    • unsupported full/vector head dims, causal qL > kL, vector GQA limits,
      and CPU execution raise with the specific violated constraint;
    • default dispatch with head_dim 192/256 full attention stays unfused and
      matches the reference;
    • head_dim=192 vector attention matches the reference.
  • Existing test_fast_sdpa.py and test_fast.py suites continue to pass.

Scope

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 (API docstring)

@apocryphx

Copy link
Copy Markdown
Contributor

Hi @hojin12312 — happy to report the M1 validation is done, and the results turned out more interesting than a simple thumbs-up. Thanks again for folding the safety and diagnostics work into your PR so cleanly; reviewing the capability-reason code while testing it on two machines was a pleasure, and I found nothing to object to — consider that portion reviewed and approved from my side.

The setup: your branch at 164fa7e23, plus a validation-only patch adding sdpa_vector 512/512 instantiations reachable through force_fused (the #3885 configuration — a test aid, not merge material). M1 Max 64 GB (applegpu_g13s) on macOS 26.5.2, with an M4 Max as control, both probing the byte-identical metallib.

1. The #3885 register cliff is still real — but it moved with the OS, and it's specialization-dependent. maxTotalThreadsPerThreadgroup for the d512 1-pass vector pipelines on the M1 Max today:

specialization bf16 fp16
no mask 1024 832
causal 832 832
bool mask 832 832
float mask 832 832

(2-pass pipelines: 1024 in every variant; M4 Max: 1024 across the board.) Back in July, every d512 variant I probed reported 832. The current Metal compiler squeezes exactly one specialization under the dispatch size; everything else still caps below the kernel's hard-coded 1024-thread launch. I did not expect the ground to shift under this experiment in three weeks, but here we are.

2. The check behaves exactly as intended, end to end (force_fused, qL=1, kL=512 → 1-pass regime on 's'-class):

  • bf16 no-mask: runs, matches the composite reference.
  • fp16 no-mask: Maximum threads per threadgroup is 832 but requested 1024 for kernel sdpa_vector_float16_t_512_512_nomask_qnt_nc_nosinks. — a loud error precisely where the pre-Check threadgroup size in the 1-pass sdpa_vector dispatch #4018 code silently returned zeros.

One reproduction note for anyone retracing this: at kL >= 1024 on 's'-class devices the dispatcher selects the 2-pass kernels (which fit everywhere), so the capped 1-pass is only reachable at shorter kL. My first run sat exactly on the wrong side of that boundary and sailed through — worth knowing before anyone concludes the hazard is gone.

3. Your force_fused tests all pass on the M1 Max — the full set of seven — plus the whole test_fast_sdpa suite on the M4 Max (25 tests, one CI-memory skip).

The takeaway I'd offer for the PR description: the hazard these checks guard is not a static device property. It shifts with the Metal compiler version and with function-constant specialization — the same kernel is safe unmasked and unsafe masked, on the same chip, today. No support-matrix allowlist can track that, which makes the unconditional dispatch-time check exactly the right shape, and your full/NAX additions close the last unchecked dispatches.

I'll close #4186 now — glad the two efforts converged, and looking forward to seeing this land. The d512 vector instantiations can then follow as the small opt-in PR discussed in #3658, with their safety story now validated on real hardware.

Comment thread mlx/fast.h Outdated
Comment thread mlx/backend/metal/scaled_dot_product_attention.cpp Outdated
Comment thread mlx/backend/metal/scaled_dot_product_attention.cpp Outdated

@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.

Thanks!

@zcbenz
zcbenz force-pushed the force-fused-sdpa branch 2 times, most recently from c7709f1 to 7830a8a Compare August 13, 2026 10:43
@zcbenz
zcbenz force-pushed the force-fused-sdpa branch 2 times, most recently from bfbf563 to f592a98 Compare August 18, 2026 10:02
@zcbenz
zcbenz merged commit 772bc8a into ml-explore:main Aug 18, 2026
55 of 56 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.

3 participants