[ROCm] Fix slowdown and Compiling.md: 1x1 convs as hipBLAS GEMM, NCHW convnet default, MIOpen naive-solver detection, ROCm 7.14 install docs

[Looong01]

Follow-up to #1234 (the unified CUDA/ROCm backend). While benchmarking the ROCm path of the merged tip on a Radeon RX 7900 XTX (gfx1100, ROCm 7.14, MIOpen 3.5.2), I found that MIOpen's NHWC convolution solver coverage is much weaker than cuDNN's, and that a few backend defaults inherited from the CUDA side cost real performance on MIOpen. This PR fixes those, and updates the ROCm install docs to the official ROCm 7.14 guide.

Problems

1. 1x1 NHWC FP16 convs are slow through MIOpen. For the transformer head convs, MIOpen picks ConvHipImplicitGemmGroupFwdXdlops (CK grouped-conv WMMA kernels) - measured ~2-4x slower per call than an equivalent GEMM, with far more dispatch overhead (thousands of extra conv dispatches per search batch).
2. FP32 1x1 NHWC convs (e.g. policy_head.conv2p) can fall all the way back to ConvDirectNaiveConvFwd - MIOpen's naive direct fallback.
3. rocmUseNHWC=Auto enabled NHWC for convnets on CDNA, but MIOpen's NCHW solver coverage (Winograd / implicit GEMM / GEMM) is its most mature path on every supported arch, including CDNA.
4. Docs: the ROCm install instructions still referenced TheRock directly; ROCm 7.14 has official packages (apt on Linux, tarball on Windows).

Changes

• rocmUse1x1Matmul (new config, mirrors cudaUse1x1Matmul, Auto = on): 1x1 NHWC convs run as hipBLAS GEMM. FP16 uses hipblasGemmEx with FP32 compute for accuracy (same reasoning as MatMulLayer); FP32 uses it too, since unlike cuDNN, MIOpen's NHWC path is weak in both precisions.
• rocmUseNHWC=Auto now resolves to NCHW for convnets on all archs. Transformer models are unchanged - they still force NHWC (the attention/RoPE/FFN implementation requires it), and their 1x1 convs are now GEMMs anyway.
• MIOpen naive-fallback detection at model load. MIOpen stamps its naive fallback solution with a sentinel time of 1000ms (registered solver id 55 as a secondary check - note the id is not public-stable API, hence two signals). If an NHWC conv's chosen solution is naive, the layer is re-expressed as NCHW with per-call transposes; a log line reports the fallback and the re-expression, so coverage gaps are visible instead of silent.
• Compiling.md: ROCm install sections (Linux + Windows) rewritten to follow https://rocm.docs.amd.com/en/latest/install/rocm.html - official apt packages on Linux (amdrocm7.14 amdrocm-core-dev7.14 ..., with the 7.2.4-or-older uninstall note and amd-smi verification) and the official tarball layout on Windows.
• Config docs for rocmUse1x1Matmul / rocmUseNHWC updated in gtp_example.cfg and analysis_example.cfg.

Scope note: all changes are inside the ROCm (KATAGO_GPU_HIP) sections or ROCm config parsing; the CUDA path is untouched.

Measurements

RX 7900 XTX (gfx1100), ROCm 7.14.0, MIOpen 3.5.2, CK 1.2.0, 19x19, FP16, benchmark at 800 visits, best visits/s across the thread sweep:

Model	master tip	this PR	delta
b10c384h6nbttflrs	2649	2835	+7.0%
b10c512h8nbt3tflrs-fson-silu-rsnh	1506	1551	+3.0%
b11c768h12nbt3tflrs-fson-silu	780	830	+6.4%
kata1-zhizi-b40c768nbt-s11272M-d5935M	354	346	neutral (within noise)

Correctness: FP16 kata-raw-nn outputs on a fixed position checked against the Eigen CPU backend as reference - whiteWin/whiteLead agree to ~1e-3 on b10c384, b10c512, b11 and the b40 convnet (e.g. b11: 0.7498 vs 0.7499).

Caveats

• The naive-fallback detection relies on MIOpen's sentinel time (1000ms), which is stable behavior, and secondarily on the registered solver id (55 for ConvDirectNaiveConvFwd in MIOpen 3.5.x), which is not guaranteed stable across MIOpen versions. If a future MIOpen renumbers it, the time sentinel still catches fallbacks.
• The NCHW re-expression adds two transposes per call; it only engages when MIOpen's NHWC options are down to naive, where even naive-plus-nothing is worse. On gfx1100 the affected layers are tiny head convs, so the transpose cost is negligible.

[lightvector]

Question: this PR deletes this code:

  if(context->useNHWCMode == enabled_t::Auto && useFP16 && isNhwcFp16PreferredArch(prop))
    useNHWC = true;

which used NHWC on some CDNA architectures. Is this deletion something that your benchmarks actually tested, or did you test RDNA only and are only guessing about this deletion? My testing found that this is a lot worse on MI300X (which I think is gfx942), on conv nets from 1.3x to 2x slower depending on the model and thread settings.

[Looong01]

Question: this PR deletes this code:

  if(context->useNHWCMode == enabled_t::Auto && useFP16 && isNhwcFp16PreferredArch(prop))
    useNHWC = true;

which used NHWC on some CDNA architectures. Is this deletion something that your benchmarks actually tested, or did you test RDNA only and are only guessing about this deletion? My testing found that this is a lot worse on MI300X (which I think is gfx942), on conv nets from 1.3x to 2x slower depending on the model and thread settings.

That deletion is based on RDNA3 (gfx1100, RX 7900 XTX) testing only. I do not have a CDNA/MI300X (gfx942) machine, so I cannot claim it is optimal there without data.

What we observed on gfx1100 is that MIOpen's NHWC path frequently falls back to the ConvDirectNaiveConv* solver, which is dramatically slower than the NCHW path. The b11c768h12 model showed a ~15% regression with NHWC auto-enabled, and the smaller conv nets were worse as well. Removing the auto-enable avoids this regression on RDNA3.

Your finding that NHWC is also worse on MI300X (1.3–2× slower on conv nets) is consistent with our observation. Do u mean this happens after this PR or before this PR? If it is before this PR, then the conservative default is the same: do not auto-enable NHWC on AMD GPUs through MIOpen. If someone later produces CDNA-specific evidence that NHWC is faster on a specific MI300X/MI325X configuration, that can be revisited with a more targeted arch guard.

[lightvector]

It happens after the PR. The specific lines I cited appear to be a deletion that cannot possibly affect gfx1100 because isNhwcFp16PreferredArch already is false on gfx1100. Why did you delete it on the basis of gfx1100 testing when such testing could not possibly exercise the effect of the deletion? And yes, deleting those lines did indeed prove to be harmful on gfx942.

[lightvector]

(Sorry, closing the PR was my laptop touchpad misfiring as I was moving the mouse)

[lightvector]

Your finding that NHWC is also worse on MI300X (1.3–2× slower on conv nets) is consistent with our observation. Do u mean this happens after this PR or before this PR?

Ah I think I understand the confusion. To be clear: NHWC is 1.3x to 2x better on MI300X. The deletion of those lines in the PR, disabling NHWC, made it 1.3x to 2x worse. And I was asking about them because their deletion only affects some of the gfx9... gpus, which wouldn't have been testable in gfx1100 benchmarks, so I was wondering why the deletion was made.

[Looong01]

@lightvector You're right — that deletion was a mistake.

I removed the CDNA auto-NHWC branch based only on RDNA3 (gfx1100) testing, where that branch already never fired, so I had no evidence to generalize it to all AMD architectures. I incorrectly treated a local RDNA3 observation as universal and it predictably hurt gfx942.

I restored it and broadened the rule: rocmUseNHWC=auto now enables NHWC for the entire gfx9* family. The implementation uses std::string_view so the prefix length is not hard-coded:

static bool isNhwcPreferredArch(const hipDeviceProp_t& prop) {
::string arch(prop.gcnArchName);
  size_t colonPos = arch.find(':');
  if(colonPos != std::string::npos)
    arch = arch.substr(0, colonPos);
  static constexpr std::string_view prefix = "gfx9";
  return arch.compare(0, prefix.size(), prefix) == 0;
}

[Looong01]

Finished

[Looong01]

@lightvector I make a little commit here: 2f1153f

Just Update CMakeLists about checking MSVC toolsets version logic for ROCm backend building on Windows. I think u can add this to our latest Update.

[lightvector]

Added.