[ROCm] Fix slowdown and Compiling.md: 1x1 convs as hipBLAS GEMM, NCHW convnet default, MIOpen naive-solver detection, ROCm 7.14 install docs - #1237
Conversation
|
Question: this PR deletes this code: 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 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. |
|
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. |
|
(Sorry, closing the PR was my laptop touchpad misfiring as I was moving the mouse) |
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
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).policy_head.conv2p) can fall all the way back toConvDirectNaiveConvFwd- MIOpen's naive direct fallback.rocmUseNHWC=Autoenabled 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.Changes
rocmUse1x1Matmul(new config, mirrorscudaUse1x1Matmul, Auto = on): 1x1 NHWC convs run as hipBLAS GEMM. FP16 useshipblasGemmExwith FP32 compute for accuracy (same reasoning asMatMulLayer); FP32 uses it too, since unlike cuDNN, MIOpen's NHWC path is weak in both precisions.rocmUseNHWC=Autonow 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.amdrocm7.14 amdrocm-core-dev7.14 ..., with the 7.2.4-or-older uninstall note andamd-smiverification) and the official tarball layout on Windows.rocmUse1x1Matmul/rocmUseNHWCupdated ingtp_example.cfgandanalysis_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,
benchmarkat 800 visits, best visits/s across the thread sweep:Correctness: FP16
kata-raw-nnoutputs 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
ConvDirectNaiveConvFwdin MIOpen 3.5.x), which is not guaranteed stable across MIOpen versions. If a future MIOpen renumbers it, the time sentinel still catches fallbacks.