Metal backend: Add gather_qmv kernel for MoE expert-indexed quantized matmul

[manuelcandales]

Adds gather_qmv Metal kernel for Mixture-of-Experts: performs per-expert
quantized matrix-vector multiply y[i] = W[expert_idx[i]] @ x[i]. Extends
the existing qmv kernels in op_linear_4bit.mm with expert
index-based pointer offsets, following the same pattern as MLX's
affine_gather_qmv_fast.

Two dispatch paths (matching op_linear_4bit.mm):

• gather_qmv_fast: optimized path for K%512==0 and N%8==0
• gather_qmv_impl: generic fallback for any K and N

Uses the same affine INT4 dequantization format as op_linear_4bit.mm
(scale * accum + sum * bias). Instantiated for 4-bit with group sizes
{32, 64, 128} and dtypes {float, bfloat16}.

Includes: Metal shader + C++ host dispatch, Python custom op definition
(metal::gather_qmv) with reference CPU impl and Meta impl, C shim dict,
fallback kernel registration, CMakeLists entry, and test module.

Authored with Claude.

[manuelcandales]

Stack from ghstack (oldest at bottom):

• Qwen 3.5 MoE: Add Metal build target #18881
• Qwen 3.5 MoE: Add --backend metal export path #18880
• Qwen 3.5 MoE: Add Metal source transformations #18879
• Metal backend: Add gated delta rule kernel for linear attention #18878
• -> Metal backend: Add gather_qmv kernel for MoE expert-indexed quantized matmul #18877
• Metal backend: Add topk fallback kernel via MPSGraph #18876
• Metal backend: Add SDPA head_dim=256 support #18875
• Metal backend: Add INT32 and BOOL dtype support #18874

[pytorch-bot]

🔗 Helpful Links

🧪 See artifacts and rendered test results at hud.pytorch.org/pr/pytorch/executorch/18877

• 📄 Preview Python docs built from this PR

Note: Links to docs will display an error until the docs builds have been completed.

❗ 1 Active SEVs

There are 1 currently active SEVs. If your PR is affected, please view them below:

• Rolling out OSDC (ARC) runners on pull workflow for PyTorch trunk commits

❌ 1 New Failure, 15 Cancelled Jobs, 8 Pending

As of commit 59f88db with merge base 5707e2a ():

NEW FAILURE - The following job has failed:

• Test CoreML Backend / test-coreml / test-backend-macos (coreml, operators) / macos-job (gh)
  RuntimeError: Command bash /Users/ec2-user/runner/_work/_temp/exec_script failed with exit code 1

CANCELLED JOBS - The following jobs were cancelled. Please retry:

• Build Presets / apple (ios-simulator) / build (gh)
• Build Presets / apple (ios) / build (gh)
• Build Presets / apple (llm) / build (gh)
• Build Presets / apple (macos) / build (gh)
• Build Presets / linux (linux, linux.arm64.2xlarge, executorch-ubuntu-22.04-gcc11-aarch64) / build (gh)
• Build Presets / linux (llm, linux.arm64.2xlarge, executorch-ubuntu-22.04-gcc11-aarch64) / build (gh)
• Build Presets / linux (pybind, linux.2xlarge, executorch-ubuntu-22.04-clang12) / build (gh)
• Build Presets / linux (pybind, linux.arm64.2xlarge, executorch-ubuntu-22.04-gcc11-aarch64) / build (gh)
• Build Presets / windows (pybind) / build (gh)
• Build Presets / zephyr (zephyr) / build (gh)
• Lint / android-java-format (gh)
• Lint / Get changed files / get-changed-files (gh)
• Lint / link-check / lint-file-size (gh)
• Lint / lintrunner-mypy (gh)
• pull (gh)

This comment was automatically generated by Dr. CI and updates every 15 minutes.