docs(rocm): document MLX_ROCM_GATHER_QMV_EXPERT_BATCHED and the other runtime MLX_ROCM_* variables #2180

Part of #1801

Problem / Background

MLX_ROCM_GATHER_QMV_EXPERT_BATCHED switches the expert-batched gather_qmm kernel, which has been on by default for bf16 and f16 affine since #2112 (#2066) and for bf16 mxfp4 since #2164 (#2106). A user who wants the per-row path back, or who is bisecting a MoE regression, can only find the variable in src/lib/mlx-cpp/patches-rocm/LOCAL_FIXES.md (items 9 and 31), two benchmark pages, a code comment and a test. It is absent from docs/environment-variables.md and from the ROCm section of docs/installation.md. The same holds for 42 other MLX_ROCM_* variables the overlay reads at runtime.

Current Behavior

• grep -rhoE 'MLX_ROCM_[A-Z0-9_]+' src/lib/mlx-cpp/patches-rocm finds 51 names. docs/environment-variables.md names 7: MLX_ROCM_ARCHITECTURES, MLX_ROCM_FFT_CACHE_SIZE, MLX_ROCM_MAX_INFLIGHT_MB, MLX_ROCM_GPU_WATCHDOG_SECS, MLX_ROCM_WMMA_QMM, MLX_ROCM_WMMA_QMM_MAX_M, MLX_ROCM_QMM_DEQUANT_GEMM, all as prose inside "Build-time variables" (docs/environment-variables.md:179-233). MLX_ROCM_HOST_STAGE is a compile definition (mlx/backend/rocm/CMakeLists.txt:458), not an environment variable, and stays undocumented. The other 43 are runtime variables with no documentation.
• docs/installation.md "Linux with AMD ROCm" (:366-574) covers only MLX_ROCM_ARCHITECTURES, MLX_ROCM_MAX_INFLIGHT_MB, MLX_ROCM_GPU_WATCHDOG_SECS and MLX_ROCM_FFT_CACHE_SIZE; its "Runtime environment variables" table (:575-591) has no ROCm rows.
• Existing text is wrong for one variable: docs/environment-variables.md:225 says the dequantize + hipBLASLt route is enabled when "MLX_ROCM_QMM_DEQUANT_GEMM not 0", but use_rocblas_dequant_path() (quantized/qmm.hip:480-492) enables it only when the variable is unset or exactly 1; any other value, including true or an empty string, disables it.

Proposed Solution

1. In docs/environment-variables.md, add a section ## ROCm backend variables directly after "Build-time variables". Move the existing runtime ROCm prose (FFT cache, in-flight bound, watchdog, WMMA QMM, WMMA QMM max M, dequant GEMM) into it unchanged except for the MLX_ROCM_QMM_DEQUANT_GEMM correction, and leave one sentence plus a link in "Build-time variables". Under it add a table with columns Variable | Values | Default | Purpose, the same as "Hardware and kernel diagnostic variables", with one row per variable below, grouped in this order: MoE gather qmv, dense qmv, dequantize + GEMM, GEMM libraries, unquantized MoE, allocator and device. State in the section intro that every path is under src/lib/mlx-cpp/patches-rocm/mlx/backend/rocm/, that most variables are read once per process and the ones marked "every call" can be changed between calls, and that "presence" variables take effect for any value, including 0.
2. In docs/installation.md, add one paragraph to the ROCm "Running" subsection naming MLX_ROCM_GATHER_QMV_EXPERT_BATCHED=0 as the way to take the per-row MoE path for A/B or bisecting, linking to the new section, and add one row to the "Runtime environment variables" table that points to it.
3. Take every value and default from the code at implementation time, not from this table, if the code has moved. fix(rocm): range-check GEMM env integers and fix a stale HIP graph comment #2152 (open) changes the GEMM integer parsers to warn on invalid input; if it merges first, document its behavior.

Scope

In scope: docs/environment-variables.md (new section, the 43 rows, the moved prose, the MLX_ROCM_QMM_DEQUANT_GEMM correction) and docs/installation.md (one paragraph, one table row).

Out of scope: code changes to any parser (#2152 covers the GEMM integers); the overlay's undocumented non-MLX_ROCM_ variables (MLX_GRAPH_*, MLX_QMV_IDOT, MLX_QMV_IDOT_GATHER, MLX_QMV_NO_WIDE, MLX_SDPA_*, MLX_HIPBLASLT_*, MLX_NO_HIPBLASLT, MLX_HSACO_CACHE_DIR, MLX_GRIDX_MULT, MLX_NO_CONCAT_SPLIT, MLX_PURE_DEBUG, MLX_COUNT_COPIES), which need their own issue.

Implementation Notes

• Reuse: the table layout of "Hardware and kernel diagnostic variables" (docs/environment-variables.md:516); the markdown conventions already in that file.
• Constraints: describe what the code does today; do not mark a variable recommended for deployment. Mark MLX_ROCM_FORCE_WARP_SIZE, MLX_ROCM_MOE_ASYNC, heuristic-only MLX_ROCM_MOE_PAD_SYNC and MLX_ROCM_GROUPED_PREFILL as debug or experimental, as their code comments do.
• Edge cases: for presence variables, say outright that =0 enables; for first-character parsers, list the accepted characters, not words.

Acceptance Criteria

• docs/environment-variables.md has a ## ROCm backend variables section with all 43 rows, each with values, default and effect matching the code, and the six runtime variables already documented in it.
• The comparison in Verification prints only MLX_ROCM_HOST_STAGE.
• The MLX_ROCM_QMM_DEQUANT_GEMM text says only unset or exactly 1 enables the route.
• docs/installation.md's ROCm section names MLX_ROCM_GATHER_QMV_EXPERT_BATCHED=0 and links to the new section, and its runtime-variable table points to it.
• Internal links resolve (anchors checked by hand or with the docs build).

Verification

comm -23 <(grep -rhoE 'MLX_ROCM_[A-Z0-9_]+' src/lib/mlx-cpp/patches-rocm/mlx src/lib/mlx-cpp/patches-rocm/CMakeLists.txt | sort -u) <(grep -oE 'MLX_ROCM_[A-Z0-9_]+' docs/environment-variables.md | sort -u)
grep -n 'MLX_ROCM_GATHER_QMV_EXPERT_BATCHED' docs/installation.md docs/environment-variables.md

A pass: the first command prints MLX_ROCM_HOST_STAGE and nothing else; the second finds the variable in both files.