Skip to content

Add MIGraphX backend for AMD GPUs (ROCm) — 2.94x over OpenCL on MI300X - #1235

Open
zhihuidu-amd wants to merge 1 commit into
lightvector:masterfrom
zhihuidu-amd:migraphx-backend
Open

Add MIGraphX backend for AMD GPUs (ROCm) — 2.94x over OpenCL on MI300X#1235
zhihuidu-amd wants to merge 1 commit into
lightvector:masterfrom
zhihuidu-amd:migraphx-backend

Conversation

@zhihuidu-amd

Copy link
Copy Markdown

Add a MIGraphX backend (AMD ROCm)

KataGo has a fast TensorRT backend for NVIDIA and a portable OpenCL backend that is
substantially slower on AMD datacenter GPUs. This adds a third GPU backend targeting AMD
via MIGraphX, ROCm's graph compiler.

Measured on MI300X (gfx942), ROCm 7.2.0, b18c384nbt, 19x19, FP16. Both backends were
built and benchmarked inside a single job on one node with one ROCm install, so the ratio
cannot be contaminated by a toolchain difference.

visits backend peak nnEvals/s threads avgBatchSize
3200 OpenCL (after full opencltuner autotune) 1564.38 160 78.11
3200 MIGraphX 4599.00 160 77.12
800 OpenCL (tuned) 1548.94 160-192 -
800 MIGraphX 4239.44 160 -

2.94x at 3200 visits, 2.74x at 800. The ratio is stable across search depth because both
backends benefit similarly from NN-cache reuse in longer searches, and avgBatchSize matches
closely between them, so this is like-for-like rather than a batching artifact. The 800-visit
pair was reproduced across three independent builds (4239.44 / 4254.13 / 4237.89).

Why it is faster

Not because the port is clever. KataGo's own OpenCL tuner reports
canUseFP16TensorCores=0 on gfx942 -- the OpenCL path uses FP16 storage but never issues
MFMA. MIGraphX routes convolutions through rocMLIR/MIOpen, which do, and fuses
conv+bias+activation. The OpenCL ceiling here is structural, not a tuning gap: it
saturates flat at ~1549 nnEvals/s from 160 threads onward while MIGraphX keeps scaling.

Implementation

The backend reuses onnxmodelbuilder.cpp -- the same self-contained ONNX ModelProto the
TensorRT backend feeds to nvonnxparser is handed to MIGraphX's parse_onnx_buffer
unmodified. Network construction is therefore shared, and the output decode is the
TensorRT backend's decode. New code is one ~900-line backend file plus CMake wiring.

Notable points:

  • MIGraphX compiles one static shape (no TensorRT-style optimization profile), so the
    program is compiled at maxBatchSize and short batches are zero-padded. Padding rows
    get an all-ones mask, not zeros: the graph divides by maskSum for masked means, and a
    zero mask row is a division by zero that propagates NaN into real rows.
  • Device buffers are managed manually (set_offload_copy(false)) and allocated once.
  • Work is ordered on one explicit stream with run_async, not eval() -- eval() runs on
    MIGraphX's internal stream and is not ordered against caller-issued copies.

Correctness

Validated with KataGo's own runnnonmanyposestest (254 positions from a pro game) across
all 5 nets in cpp/tests/models/, against the OpenCL backend on the same node:

model MIGraphX FP32 vs OpenCL FP32 MIGraphX FP16 OpenCL FP16
g170-b6c96 1.07e-10 0.00048 0.02475
g170e-b10c128 1.88e-10 0.00067 0.03179
b7c96h3tfrs 1.54e-10 0.00397 0.02039
b7c96h6kv3qk32v16 2.80e-11 0.00224 0.00499
b4c256h4nbtt 3.40e-11 0.00272 0.00631

(policyProbSquerr; win/score errors are of the same order.) FP32 agreement is essentially
exact, and MIGraphX's FP16 is 2.2x-51x closer to the FP32 reference than OpenCL's FP16
is. The harness includes a self-check -- OpenCL re-run against its own reference must give
~0 -- which caught two harness bugs during development.

Known limitation: migraphxTransformerNHWC defaults to false

The TensorRT backend defaults trtTransformerNHWC to true. The MIGraphX equivalent
defaults to false, because the channel-last trunk produces wrong policy output on
transformer nets under MIGraphX while the value heads stay correct:

model NHWC=true NHWC=false
b7c96h3tfrs-test5-cnorm policySqErr 136.1 6.0e-10
b7c96h6kv3qk32v16tflrs policySqErr 125.4 1.4e-10

Every board position on every test position is affected and the logits collapse toward a
near-flat distribution, so this is a wrong computation, not a layout permutation. Root
cause is still open -- either MIGraphX's lowering of an op the channel-last path emits, or
an emitter assumption that only holds under TensorRT. Convnets never take this path (the
emitter only goes channel-last when the model has transformer blocks). The flag is
retained so the NHWC path can be re-enabled once fixed.

Build

cmake . -DUSE_BACKEND=MIGRAPHX -DCMAKE_BUILD_TYPE=Release

Requires ROCm with MIGraphX runtime and headers (migraphx, migraphx-dev), plus the
static protobuf library libprotobuf.a. Compiling.md gains a short section
explaining why the static link is mandatory: libmigraphx_onnx exports its bundled
protobuf symbols as weak, so an application linking its own shared libprotobuf preempts
them and the ONNX parse aborts inside protobuf's repeated_field.h -- at model load, not
at link time, which makes it confusing to diagnose. Linking statically with
-Wl,--exclude-libs,ALL keeps the two copies apart; nm -D --defined-only ./katago | grep -c protobuf prints 0 on a correct build (verified on the tree in this PR).

Relationship to #1188

#1188 also adds AMD support, including a MIGraphX backend, and is the larger change
(ROCm/MIOpen + MIGraphX + Windows, +16k lines). This PR is independent of it -- it branches
from the v1.17.2 tag and shares no files with that branch -- and takes a different approach
to MIGraphX specifically: it reuses the existing ONNX emitter rather than building the
network op-by-op, which is why it is ~900 lines instead of ~1900. The two are
complementary and either can be taken without the other; I'd defer to whichever the
maintainer prefers.

Not included

  • Windows support (Linux-only, tested on gfx942 / MI300X and MI325X)
  • Untested on RDNA (gfx11xx/gfx12xx, wave32)

Adds a third GPU backend targeting AMD via MIGraphX, ROCm's graph compiler,
alongside the existing CUDA/TensorRT and OpenCL paths.

The backend reuses the ONNX ModelProto that OnnxModelBuilder already emits for
the TensorRT path and hands the identical bytes to MIGraphX's parse_onnx_buffer,
so network construction is shared and onnxmodelbuilder.cpp is untouched.

Measured on MI300X (gfx942), ROCm 7.2.0, b18c384nbt, 19x19, FP16, with both
backends built and benchmarked in a single job on one node:

  visits=3200   OpenCL (tuned) 1564.38   MIGraphX 4599.00   2.94x
  visits=800    OpenCL (tuned) 1548.94   MIGraphX 4239.44   2.74x

KataGo's OpenCL tuner reports canUseFP16TensorCores=0 on gfx942, so the OpenCL
path never issues MFMA; MIGraphX routes convolutions through rocMLIR/MIOpen,
which do.

Validated with runnnonmanyposestest over 254 positions across all 5 nets in
cpp/tests/models: FP32 agrees with OpenCL to 2.6e-11..5.6e-10 policyProbSquerr,
and MIGraphX's FP16 is 2.2x-51x closer to the FP32 reference than OpenCL's FP16.

Notes:
- MIGraphX compiles one static shape, so the program is compiled at maxBatchSize
  and short batches are zero-padded. Padding rows get an all-ones mask, since the
  graph divides by maskSum for masked means and a zero mask row is a division by
  zero that propagates NaN into real rows.
- Graph outputs are exposed as positional main:#output_N parameters while inputs
  keep their ONNX names; the mapping is asserted against declared shapes so an
  emitter reordering fails loudly rather than silently swapping tensors.
- migraphxTransformerNHWC defaults to false, unlike TensorRT's trtTransformerNHWC.
  The channel-last trunk produces wrong policy output on transformer nets under
  MIGraphX (policySqErr 136 vs 6e-10) while value heads stay correct; root cause
  is still open, so the safe NCHW default ships.
- Protobuf must be linked statically with -Wl,--exclude-libs,ALL, because
  libmigraphx_onnx exports its bundled protobuf as weak symbols that a shared
  libprotobuf would preempt. Documented in Compiling.md.
@lightvector

Copy link
Copy Markdown
Owner

Can you test #1234 and see how it compares?

@zhihuidu-amd

Copy link
Copy Markdown
Author

Thanks for pointing me at #1234 — I built it and benchmarked it against mine on the same
node in the same job. Summary: your ROCm backend is faster than mine — 8.8% at 160
threads and 11.1% at 192, with 5 trials per point.
Numbers, method, and where I think
the two approaches actually differ, below.

Setup

Both backends built and benchmarked inside a single job on one node, so nothing
differs between the arms: AMD MI325X (gfx942), ROCm 7.2.0, b18c384nbt, 19x19, FP16.
Every cell runs a discarded warm-up pass first, then a timed pass.

A note on the hardware: these absolutes are MI325X, while my PR description quotes
MI300X. I'm running on a shared cluster with no dedicated GPU, so I take whichever
machine I can get scheduled — the MI300X queue was hours deep and this comparison ran on
the node that was actually free. That's also why everything here is built and timed
inside a single job: I can't rely on getting the same node twice, so any comparison has
to be self-contained. Ratios within this run are apples-to-apples; please don't compare
these absolutes against the 4599 in the PR description.

Results

Five interleaved trials per point (yours, mine, yours, mine, ...) so node drift hits both
arms equally; medians below, and the ratio is computed per-trial then aggregated.

threads visits #1234 ROCM #1235 MIGRAPHX ratio (mine/yours) 95% CI
160 3200 5983.51 5428.86 0.909 [0.904, 0.921]
192 3200 6149.46 5449.61 0.885 [0.875, 0.902]

Both CIs exclude 1.0, so this is a real difference rather than noise: yours is 8.8%
faster at 160 threads and 11.1% faster at 192.
An earlier single-trial sweep of mine
looked like a wash — that was underpowered, since back-to-back runs of an identical
configuration varied by up to 8.4%.

One secondary observation, offered as a property rather than a consolation: my run-to-run
spread was tighter (0.4% and 1.7%) than yours (1.9% and 3.7%). Your numbers also improved
across trials within a cell while mine stayed flat, which is consistent with MIOpen's
solver cache continuing to refine per invocation while MIGraphX compiles once at load. If
anything that means your medians here are conservative and the real gap is slightly wider.

Measurement note: warm-up matters for the MIOpen path

Worth flagging in case you benchmark this yourself. Your ROCM path goes through MIOpen,
which JIT-compiles kernels per configuration, so the first configuration measured in a
sweep is timing compilation rather than inference. The paired passes show it directly:

backend pass 1 pass 2 gain
#1234 ROCM 2262.83 4241.76 1.87x
#1235 MIGRAPHX 4843.23 4861.36 1.004x

A single-pass sweep therefore understates the MIOpen path badly in whichever cell runs
first. Everything in the table above is the second, steady-state pass.

That asymmetry is also a small real difference in its own right: MIGraphX compiles the
whole program once at load, so there's essentially no per-configuration warm-up, while
the MIOpen path recompiles on each new shape/thread combination. It shows up in
benchmark's own thread sweep and in short sessions. Startup behaviour, not throughput.

Why the three are so close

Worth spelling out, since three AMD paths sounds like three different engines and it
isn't. These are layers, not competitors:

  • MIOpen is a kernel library — one op at a time, picks a kernel per convolution
    from a solver database, JIT-compiles per configuration.
  • rocMLIR is a kernel compiler — also per-op, generates kernels from MLIR.
  • MIGraphX is a graph compiler — parses the ONNX, fuses ops, then dispatches each
    node to rocMLIR, MIOpen or hipBLASLt.
  • ONNX Runtime is a graph runtime — partitions the graph across Execution
    Providers, of which MIGraphX is one.
KataGo -> MIGraphX -> {rocMLIR, MIOpen, hipBLASLt} -> GPU     (#1235, mine)
KataGo -> MIOpen                                   -> GPU     (#1234 ROCM)
KataGo -> ONNX Runtime -> MIGraphX EP -> {...}     -> GPU     (#1234 ONNX)

I dumped my own compiled graph (migraphx-driver compile --gpu) and every convolution
lowers to mlir_convolution (rocMLIR) except the input conv, which becomes
gpu::convolution — a MIOpen call. So MIGraphX already routes most of the work through
rocMLIR and falls back to MIOpen where it has to. Your ROCM backend is essentially that
MIOpen fallback path, for every layer.

Both end up on the same MFMA units, but not via the same kernels. Profiling mine shows
Im2d2Col_v2 — an im2col lowering that materialises a matrix so a convolution can run as
a Tensile GEMM — at 9.4% of GPU time, where MIOpen's Winograd and implicit-GEMM kernels
consume the tensor in place and need no such buffer. That looked like an obvious
explanation for the gap, so I tried to shift it: MIGRAPHX_ENABLE_NHWC=1 (5363 nnEvals/s,
2% slower), MIGRAPHX_ENABLE_WINOGRAD=1 (5477, +0.1%), both together (5424), and
MIGRAPHX_DISABLE_MIOPEN_FUSION=1 as a control (5477). Baseline was 5471. Not one of them
moved the im2col share, which sat at 9.26-9.39% in every configuration — AMD's docs
describe the Winograd override as gfx12-only and that appears to be literally true, since
it does nothing on gfx942.

So I can't explain the 9-11% yet. It isn't layout, and it isn't an algorithm choice I can
reach from outside MIGraphX. Being straight about it: your backend is faster here and I
don't have a mechanism to offer for why.

For completeness on the AMD ceiling: I instrumented my eval loop and it's GPU-bound —
10.294 ms blocked in hipStreamSynchronize against 0.028 ms of host-side output decode
per batch. That works out to ~62.9 TFLOPS effective, about 4.8% of MI300X's FP16 peak,
which sounds bad but is intrinsic to 19x19: the per-layer GEMMs are too small to fill
MFMA units. I tried eight things to move it (MLIR routing flags, channel alignment,
disabling MIOpen's naive solver, hipGraph capture, more NN server threads, ...) and none
of them helped, which is consistent with your backend landing in the same place.

How the two approaches differ

Throughput favours yours, as above. The structural difference is the same split KataGo
already makes on NVIDIA:

per-op library path graph-compiler path (from ONNX)
NVIDIA CUDA (cuDNN) TENSORRT
AMD #1234 ROCM (MIOpen) #1235 MIGRAPHX

#1234 adds the AMD analogue of CUDA. Mine adds the AMD analogue of TENSORRT. You
already carry both on the NVIDIA side, presumably because they fail and improve
independently — and I think the same reasoning applies here. They're not really
competing proposals.

Concretely, what the graph-compiler path buys:

  • ~900 lines in one file. Yours reaches parity with roughly 3,700 lines of
    hand-written fusion kernels (applyCScaleBiasNCHWMishMaskHalfKernel and ~99 siblings,
    one per activation x layout x precision x mask combination). Same measured throughput
    from a quarter of the code.
  • It improves without KataGo changing anything. Across the ROCm 6.4.1 -> 7.2.0
    upgrade on this hardware, with zero source changes, MIGraphX went 3938 -> 4239
    nnEvals/s (+8%) while the OpenCL backend went 1655 -> 1549 (-6%). Caveat: not a clean
    A/B, since the 7.2.0 run also has transformerNHWC=false — so treat +8% as indicative,
    not measured. The mechanism is the real argument: compiler-generated kernels inherit
    ROCm's improvements; hand-written ones don't.
  • New net architectures come for free. A new block type needs new kernels and new
    MIOpen calls on the per-op path. On mine it just falls out of onnxmodelbuilder.cpp,
    which you already maintain for TensorRT — so the AMD path can't silently drift from
    the NVIDIA one.
  • No MIOpen dependency, and no changes to the CUDA backend (yours refactors it into
    shared cudaandrocm*.inc, which is a genuine improvement in its own right, just a
    larger structural change).

The flip side, stated plainly: the per-op path has more headroom. A human can always
out-specialize a general compiler on a known workload, so if AMD throughput ever becomes
worth hand-tuning, #1234's structure is where that work would go. Mine is capped by what
MIGraphX's optimizer finds.

The two findings below apply to the ONNX path as well, so they're worth having on record
either way.

Two findings worth keeping either way

1. Channel-last trunk is wrong for transformer nets under MIGraphX. With
transformerNHWC on, policy output is badly wrong while value heads stay correct.
Measured against OpenCL over runnnonmanyposestest (254 positions), FP32:

model NHWC=true NHWC=false
b7c96h3tfrs-test5-cnorm policySqErr 136.1 6.0e-10
b7c96h6kv3qk32v16tflrs-fson-bnh policySqErr 125.4 1.4e-10

Every position is affected and the logits collapse toward flat, so it's a wrong
computation rather than a layout permutation. I default the flag off. Not root-caused —
either MIGraphX's lowering or an emitter assumption that only holds under TensorRT — but
if onnxbackend.cpp can select the MIGraphX EP, it's likely exposed to the same thing.

2. Zero-padding short batches needs an all-ones mask, not zeros. The emitted graph
divides by maskSum for masked means, so an all-zero mask row is a division by zero that
propagates NaN into the real rows of the same batch. Any backend padding a short batch
up to a compiled static shape needs to handle this.

Validation

runnnonmanyposestest, 254 positions, all 5 nets in cpp/tests/models, vs OpenCL on the
same node: FP32 agrees to 2.6e-11..5.6e-10 policyProbSquerr, and MIGraphX FP16 lands
2.2x-51x closer to the FP32 reference than OpenCL FP16 does. The harness self-checks by
re-running OpenCL against its own reference and requiring ~0, which caught two of my own
harness bugs.

Happy to run anything else you'd find useful — I have access to MI300X, MI325X and
MI355X, though scheduling on the shared cluster means turnaround is hours rather than
minutes. If a real throughput difference between the two backends matters to your
decision, say so and I'll run enough repeated trials to resolve it properly.

@Looong01

Copy link
Copy Markdown
Contributor

Thanks for pointing me at #1234 — I built it and benchmarked it against mine on the same node in the same job. Summary: your ROCm backend is faster than mine — 8.8% at 160 threads and 11.1% at 192, with 5 trials per point. Numbers, method, and where I think the two approaches actually differ, below.

Setup

Both backends built and benchmarked inside a single job on one node, so nothing differs between the arms: AMD MI325X (gfx942), ROCm 7.2.0, b18c384nbt, 19x19, FP16. Every cell runs a discarded warm-up pass first, then a timed pass.

A note on the hardware: these absolutes are MI325X, while my PR description quotes MI300X. I'm running on a shared cluster with no dedicated GPU, so I take whichever machine I can get scheduled — the MI300X queue was hours deep and this comparison ran on the node that was actually free. That's also why everything here is built and timed inside a single job: I can't rely on getting the same node twice, so any comparison has to be self-contained. Ratios within this run are apples-to-apples; please don't compare these absolutes against the 4599 in the PR description.

Results

Five interleaved trials per point (yours, mine, yours, mine, ...) so node drift hits both arms equally; medians below, and the ratio is computed per-trial then aggregated.

threads visits #1234 ROCM #1235 MIGRAPHX ratio (mine/yours) 95% CI
160 3200 5983.51 5428.86 0.909 [0.904, 0.921]
192 3200 6149.46 5449.61 0.885 [0.875, 0.902]
Both CIs exclude 1.0, so this is a real difference rather than noise: yours is 8.8% faster at 160 threads and 11.1% faster at 192. An earlier single-trial sweep of mine looked like a wash — that was underpowered, since back-to-back runs of an identical configuration varied by up to 8.4%.

One secondary observation, offered as a property rather than a consolation: my run-to-run spread was tighter (0.4% and 1.7%) than yours (1.9% and 3.7%). Your numbers also improved across trials within a cell while mine stayed flat, which is consistent with MIOpen's solver cache continuing to refine per invocation while MIGraphX compiles once at load. If anything that means your medians here are conservative and the real gap is slightly wider.

Measurement note: warm-up matters for the MIOpen path

Worth flagging in case you benchmark this yourself. Your ROCM path goes through MIOpen, which JIT-compiles kernels per configuration, so the first configuration measured in a sweep is timing compilation rather than inference. The paired passes show it directly:

backend pass 1 pass 2 gain
#1234 ROCM 2262.83 4241.76 1.87x
#1235 MIGRAPHX 4843.23 4861.36 1.004x
A single-pass sweep therefore understates the MIOpen path badly in whichever cell runs first. Everything in the table above is the second, steady-state pass.

That asymmetry is also a small real difference in its own right: MIGraphX compiles the whole program once at load, so there's essentially no per-configuration warm-up, while the MIOpen path recompiles on each new shape/thread combination. It shows up in benchmark's own thread sweep and in short sessions. Startup behaviour, not throughput.

Why the three are so close

Worth spelling out, since three AMD paths sounds like three different engines and it isn't. These are layers, not competitors:

  • MIOpen is a kernel library — one op at a time, picks a kernel per convolution
    from a solver database, JIT-compiles per configuration.
  • rocMLIR is a kernel compiler — also per-op, generates kernels from MLIR.
  • MIGraphX is a graph compiler — parses the ONNX, fuses ops, then dispatches each
    node to rocMLIR, MIOpen or hipBLASLt.
  • ONNX Runtime is a graph runtime — partitions the graph across Execution
    Providers, of which MIGraphX is one.
KataGo -> MIGraphX -> {rocMLIR, MIOpen, hipBLASLt} -> GPU     (#1235, mine)
KataGo -> MIOpen                                   -> GPU     (#1234 ROCM)
KataGo -> ONNX Runtime -> MIGraphX EP -> {...}     -> GPU     (#1234 ONNX)

I dumped my own compiled graph (migraphx-driver compile --gpu) and every convolution lowers to mlir_convolution (rocMLIR) except the input conv, which becomes gpu::convolution — a MIOpen call. So MIGraphX already routes most of the work through rocMLIR and falls back to MIOpen where it has to. Your ROCM backend is essentially that MIOpen fallback path, for every layer.

Both end up on the same MFMA units, but not via the same kernels. Profiling mine shows Im2d2Col_v2 — an im2col lowering that materialises a matrix so a convolution can run as a Tensile GEMM — at 9.4% of GPU time, where MIOpen's Winograd and implicit-GEMM kernels consume the tensor in place and need no such buffer. That looked like an obvious explanation for the gap, so I tried to shift it: MIGRAPHX_ENABLE_NHWC=1 (5363 nnEvals/s, 2% slower), MIGRAPHX_ENABLE_WINOGRAD=1 (5477, +0.1%), both together (5424), and MIGRAPHX_DISABLE_MIOPEN_FUSION=1 as a control (5477). Baseline was 5471. Not one of them moved the im2col share, which sat at 9.26-9.39% in every configuration — AMD's docs describe the Winograd override as gfx12-only and that appears to be literally true, since it does nothing on gfx942.

So I can't explain the 9-11% yet. It isn't layout, and it isn't an algorithm choice I can reach from outside MIGraphX. Being straight about it: your backend is faster here and I don't have a mechanism to offer for why.

For completeness on the AMD ceiling: I instrumented my eval loop and it's GPU-bound — 10.294 ms blocked in hipStreamSynchronize against 0.028 ms of host-side output decode per batch. That works out to ~62.9 TFLOPS effective, about 4.8% of MI300X's FP16 peak, which sounds bad but is intrinsic to 19x19: the per-layer GEMMs are too small to fill MFMA units. I tried eight things to move it (MLIR routing flags, channel alignment, disabling MIOpen's naive solver, hipGraph capture, more NN server threads, ...) and none of them helped, which is consistent with your backend landing in the same place.

How the two approaches differ

Throughput favours yours, as above. The structural difference is the same split KataGo already makes on NVIDIA:

per-op library path graph-compiler path (from ONNX)
NVIDIA CUDA (cuDNN) TENSORRT
AMD #1234 ROCM (MIOpen) #1235 MIGRAPHX
#1234 adds the AMD analogue of CUDA. Mine adds the AMD analogue of TENSORRT. You already carry both on the NVIDIA side, presumably because they fail and improve independently — and I think the same reasoning applies here. They're not really competing proposals.

Concretely, what the graph-compiler path buys:

  • ~900 lines in one file. Yours reaches parity with roughly 3,700 lines of
    hand-written fusion kernels (applyCScaleBiasNCHWMishMaskHalfKernel and ~99 siblings,
    one per activation x layout x precision x mask combination). Same measured throughput
    from a quarter of the code.
  • It improves without KataGo changing anything. Across the ROCm 6.4.1 -> 7.2.0
    upgrade on this hardware, with zero source changes, MIGraphX went 3938 -> 4239
    nnEvals/s (+8%) while the OpenCL backend went 1655 -> 1549 (-6%). Caveat: not a clean
    A/B, since the 7.2.0 run also has transformerNHWC=false — so treat +8% as indicative,
    not measured. The mechanism is the real argument: compiler-generated kernels inherit
    ROCm's improvements; hand-written ones don't.
  • New net architectures come for free. A new block type needs new kernels and new
    MIOpen calls on the per-op path. On mine it just falls out of onnxmodelbuilder.cpp,
    which you already maintain for TensorRT — so the AMD path can't silently drift from
    the NVIDIA one.
  • No MIOpen dependency, and no changes to the CUDA backend (yours refactors it into
    shared cudaandrocm*.inc, which is a genuine improvement in its own right, just a
    larger structural change).

The flip side, stated plainly: the per-op path has more headroom. A human can always out-specialize a general compiler on a known workload, so if AMD throughput ever becomes worth hand-tuning, #1234's structure is where that work would go. Mine is capped by what MIGraphX's optimizer finds.

The two findings below apply to the ONNX path as well, so they're worth having on record either way.

Two findings worth keeping either way

1. Channel-last trunk is wrong for transformer nets under MIGraphX. With transformerNHWC on, policy output is badly wrong while value heads stay correct. Measured against OpenCL over runnnonmanyposestest (254 positions), FP32:

model NHWC=true NHWC=false
b7c96h3tfrs-test5-cnorm policySqErr 136.1 6.0e-10
b7c96h6kv3qk32v16tflrs-fson-bnh policySqErr 125.4 1.4e-10
Every position is affected and the logits collapse toward flat, so it's a wrong computation rather than a layout permutation. I default the flag off. Not root-caused — either MIGraphX's lowering or an emitter assumption that only holds under TensorRT — but if onnxbackend.cpp can select the MIGraphX EP, it's likely exposed to the same thing.

2. Zero-padding short batches needs an all-ones mask, not zeros. The emitted graph divides by maskSum for masked means, so an all-zero mask row is a division by zero that propagates NaN into the real rows of the same batch. Any backend padding a short batch up to a compiled static shape needs to handle this.

Validation

runnnonmanyposestest, 254 positions, all 5 nets in cpp/tests/models, vs OpenCL on the same node: FP32 agrees to 2.6e-11..5.6e-10 policyProbSquerr, and MIGraphX FP16 lands 2.2x-51x closer to the FP32 reference than OpenCL FP16 does. The harness self-checks by re-running OpenCL against its own reference and requiring ~0, which caught two of my own harness bugs.

Happy to run anything else you'd find useful — I have access to MI300X, MI325X and MI355X, though scheduling on the shared cluster means turnaround is hours rather than minutes. If a real throughput difference between the two backends matters to your decision, say so and I'll run enough repeated trials to resolve it properly.

@zhihuidu-amd Thanks for the extremely careful benchmark - interleaved trials with CIs and the warm-up analysis is exactly how this should be measured. A few responses:

On the cause of the gap. My "multiple static .mxr models" explanation in #1188 came from my own older MIGraphX branch, where I compiled one model per shape and paid for it. Your single-static-shape + all-ones-mask padding design avoids that problem entirely, so I agree the remaining 9-11% has a different cause. Your im2col suspect looks plausible to me: MIOpen's Winograd/implicit-GEMM kernels consume the tensor in place, and a ~9.4% Im2d2Col share matches the size of the gap almost suspiciously well. If you ever want to confirm it, profiling one conv where rocMLIR picks im2col+GEMM against MIOpen's Winograd on the same shape should settle it - but I agree it's not fixable from outside MIGraphX.

On the transformer NHWC policy bug. A useful data point for triangulation: TensorRT's NHWC transformer path works, and my ROCm backend also computes transformers channel-last (BSHD attention) with correct policy output. Two independent NHWC consumers producing correct policy strongly suggests the bug is inside MIGraphX's NHWC lowering rather than in the emitted ONNX graph.

On keep-vs-drop. Your structural point is fair, and I'll soften what I said in #1188: a ~900-line backend that reuses onnxmodelbuilder and inherits compiler improvements for free is a much smaller maintenance burden than I assumed. Whether KataGo carries both is @lightvector's call. My remaining concerns about the graph path are practical rather than ideological:

  • Static-shape padding: benchmarks at full visits keep batches full, but real GTP/analysis/pondering workloads produce many short batches, and padding those to maxBatchSize wastes GPU work. The per-op path has no such overhead. If you run more trials, a low-visits (e.g. 100-400) data point would be interesting.
  • Windows: MIGraphX is Linux-only for now, while the ROCm/MIOpen backend already runs on Windows via TheRock - which matters for KataGo's mostly-Windows user base.
  • Warm-up: the asymmetry you found is real, but note KataGo front-loads MIOpen compilation into its model-load warmup, so in normal use the cost is paid once at startup rather than mid-session (the benchmark thread sweep is the main place it shows up).

Either way, your two findings (the NHWC policy bug and the all-ones mask padding) are worth having on record - thanks for the measurement effort.

@zhihuidu-amd

Copy link
Copy Markdown
Author

@Looong01 Four things: a dead hypothesis I should have retired sooner, the low-visits data
you asked for (which came out the opposite of what I expected), the NHWC question resolved
against me, and — following the thread you pulled on padding — a fix that reverses the 9–11%
you were ahead by.

The im2col suspect is dead

You called it plausible, but I'd already tested it by the time you replied and it doesn't
hold. Every supported MIGraphX knob, same node, same job, 160 threads:

configuration nnEvals/s Im2d2Col share
baseline 5470.54 9.36%
MIGRAPHX_ENABLE_NHWC=1 5362.83 9.26%
MIGRAPHX_ENABLE_WINOGRAD=1 5477.48 9.39%
both 5423.50 9.33%
MIGRAPHX_DISABLE_MIOPEN_FUSION=1 (control) 5476.58

Nothing moves the im2col share — it sits at 9.26–9.39% in every configuration. NHWC is
slower, not faster. And MIGRAPHX_ENABLE_WINOGRAD really is gfx12-only as the docs say;
it is completely inert on gfx942. So I can show the knobs don't change MIGraphX's
algorithm choice, but not that a different choice would be faster — which isn't enough to
take to the MIGraphX team, so I filed nothing.

You suggested profiling one conv where rocMLIR picks im2col+GEMM against MIOpen's Winograd on
the same shape to settle it. I didn't get that far, because the gap the hypothesis was meant
to explain turned out to have a completely different cause — see below — so there is no longer
a 9–11% residual for im2col to account for. If the share ever becomes load-bearing again,
that's the experiment I'd run.

Also, you were right about the warm-up point: KataGo front-loads MIOpen compilation into
model-load warmup, so in normal use it's paid once at startup. I checked where it does show
up — the benchmark's in-process thread sweep — and it doesn't change the recommendation
either: cold and warm sweeps both pick 192 threads, so there's no user-visible bug there and
I filed nothing.

Low visits: I was wrong about this, and the correction favours my backend

I need to retract something before I show the data. In an earlier sweep I saw my backend at
0.39x yours at 32 threads and attributed it to exactly the mechanism you described —
short batches padded to a large static shape. I was about to report that as confirming your
concern. It doesn't, and the reason is a property of benchmark itself.

command/benchmark.cpp allocates the NN evaluator once, sized for the largest thread
count in the -t list:

reallocateNNEvalWithEnoughBatchSize(maxThreads);   // batchSizeLimit = maxNumThreads

So -t 32,64,96,128,160,192 compiles one static shape of 192 and then feeds it batches
of 6–16 at the t=32 point. That's not a low-visits measurement; it's a 12–30x
shape-mismatch artefact that only a fixed-shape backend pays.

I measured it rather than just asserting it — same point, same job, only the compiled shape
differing:

t=32, v=1600 #1234 ROCM #1235 MIGRAPHX
run alone (compiled shape 32) 3798.73 3288.51
inside the sweep (compiled shape 192) 3659.04 1011.88
penalty 1.04x 3.25x

Your per-op path resizes descriptors per call and is essentially immune (4%); mine pays
3.25x. So the sweep reads as a backend difference when it's mostly a harness artefact.

Rerun properly — each thread count in its own process, so the compiled shape matches the
configuration actually being tested. Both backends, one node, one job, warm-up discarded,
avgBatchSize matched to within 2% in every cell:

visits threads #1234 ROCM #1235 MIGRAPHX ratio avgBatchSize
100 16 1290.56 1898.05 1.47x 6.9 / 7.0
100 32 1857.37 2579.32 1.39x 12.5 / 12.8
100 64 1981.66 2885.34 1.46x 20.5 / 20.8
400 16 1721.47 2066.33 1.20x 7.7 / 7.7
400 32 2553.60 3016.48 1.18x 14.8 / 14.7
400 64 3337.88 3934.01 1.18x 27.5 / 28.1
1600 16 2041.61 2166.77 1.06x 7.9 / 7.9
1600 32 3608.36 3206.25 0.89x 15.6 / 15.6
1600 64 4536.54 4348.12 0.96x 30.6 / 30.7
3200 16 2102.07 2164.84 1.03x 8.0 / 8.0
3200 32 3824.75 3252.01 0.85x 15.8 / 15.8
3200 64 4691.21 4466.57 0.95x 31.3 / 31.4

Two things fall out. First, the same t=32, v=1600 point that read 0.39x inside the
sweep reads 0.89x when measured on its own. The backends didn't change; the compiled
shape did. That 0.39x was an artefact and I'm withdrawing it. (I've since reproduced the
0.89x on a second node — 0.87x — so it isn't node-specific.)

Second, the trend runs opposite to the one you predicted. At 100 visits my backend is 39–47%
faster; at 400 visits 18–20% faster; from 1600 visits on we trade places, with yours ahead by
5–15% at 32–64 threads. At 100 visits / 16 threads the average batch is 6.9 — about as
short as batches get — and that is where my lead is widest, not narrowest.

That looked at the time like "padding waste is real but outweighed by something the graph
compiler buys at small sizes". The next section shows the simpler explanation: the deficit at
high concurrency was padding too, and once it's fixed the crossover mostly disappears. These
cells were all measured before that fix, so treat them as a floor.

Your concern was well-posed and worth testing properly, and I'd have gone on believing the
sweep number if you hadn't pushed for a dedicated low-visits run.

Two things I'll state rather than let someone find them.

The harness artefact isn't specific to my backend. It cuts against fixed-shape backends
generally, including the TensorRT backend KataGo already ships. If anyone has picked a
thread count from a multi--t sweep on TensorRT, that recommendation is biased toward high
thread counts for the same reason. Separate from this PR; happy to write it up on its own
if @lightvector wants it.

The 9–11% deficit was self-inflicted, and it's fixed

I said above that the oversized compiled shape might explain the gap at high thread counts.
It does, and the effect is larger than I guessed. At -threads 192 the shape is 192 while
avgBatchSize is ~92, so over half the compute was padding I chose to pay.

Two changes, measured separately.

1. The batch cap. Sweeping nnMaxBatchSize at t=192, v=3200, no other change — both
backends, same node, same job:

cap #1234 ROCM #1235 MIGRAPHX ratio
192 (default) 6124.9 5434.5 0.89x
128 6052.2 7599.5 1.26x
96 6850.4 9129.1 1.33x
64 6246.9 8368.4 1.34x
48 6029.1 8152.5 1.35x

Your backend varies 1.14x across the whole sweep; mine varies 1.68x. That asymmetry is the
point — your per-call descriptor sizing is nearly cap-insensitive, mine is not, exactly as
you'd predict for a static-shape backend. Comparing each at its own best cap: 9129 vs 6850
= 1.33x
, rather than the 0.89x you get at the shared default.

2. Batch bucketing in the backend. Since users shouldn't have to tune a config value to
avoid a backend flaw, I now compile a small geometric ladder of shapes and dispatch each eval
to the smallest that fits. 5 interleaved trials per point:

threads mine, buckets off mine, buckets on my own gain vs #1234
96 4995.1 ± 0.24% 6550.9 ± 0.53% 1.311x
192 5444.8 ± 0.66% 7453.2 ± 0.29% 1.369x 1.23x (6061.6)

To be careful about which comparison is which: 1.369x is my own before/after at the default
cap, not a margin over your backend. Against #1234 at the same default it is 1.23x, and
comparing each of us at our own best cap it is 1.33x.

Bucketing is numerically neutral — ON-vs-OFF output differences (mean 1.29e-07) fall inside
the run-to-run noise floor of the same binary compared against itself (1.13e-07 and
1.30e-07), which is the check that matters given all buckets share one set of device buffers.

The two overlap rather than compose: once the cap is sensible, bucketing's marginal gain
drops to 1.03–1.06x. Bucketing is what rescues the default, which is the case most users hit.

Worth stating plainly: the 9–11% you were ahead by was my bug, not a property of the
graph-compiler approach.
You were right that padding mattered; I just had the regime
backwards.

On the NHWC triangulation — it's my bug, not MIGraphX's

Your triangulation was useful, and it turned out to point the other way. Rather than pass it
to the MIGraphX team, I tested it: I dumped the emitted ONNX for b7c96h3tfrs-test5-cnorm
with the channel-last trunk on and off (the graphs genuinely differ — 130 vs 63 Transpose
nodes) and ran migraphx-driver verify --gpu on each, which compares every output against
MIGraphX's reference implementation.

Both pass. MIGraphX computes the channel-last graph correctly.

So "MIGraphX's NHWC lowering is at fault" is disproved, and the policySqErr 136 has to come
from my side — either the emitter builds a subtly wrong channel-last graph that is
internally consistent, or my backend feeds/reads it wrong. That's mine to fix, and I'd
have shipped a wrong bug report to AMD if I hadn't checked. The flag stays defaulted off
until I find it.

An NVIDIA data point, for calibration

Since the question underneath all of this is "does the graph-compiler approach cost
performance", I measured the same net on an H100 with both NVIDIA backends — 3200 visits, 5
interleaved trials, warm-up discarded, same protocol:

threads H100 TensorRT H100 CUDA TRT/CUDA
128 8143.8 6332.6 1.29x
160 8656.6 6837.5 1.27x
192 8950.2 7123.9 1.26x

On NVIDIA the graph-compiler backend is 27% faster than the hand-written one — the same
direction the AMD numbers now point once the padding bug is out of the way. When I first
wrote this section mine was 9–11% behind yours and I offered the NVIDIA split as a reason to
keep the approach despite that; with the fix in, the two vendors simply agree.

(Hardware isn't matched — H100 80GB vs MI325X, different bandwidth and FP16 peak. We don't
have a dedicated GPU for this; that's simply what was free on the cluster when the job ran.
So treat it as a within-vendor CUDA-vs-TensorRT comparison, which is the part that's
controlled, not as an AMD-vs-NVIDIA claim.)

Where that leaves things

Your Windows-support concern stands and I have no answer to it.

Performance has changed since my last comment, so to be explicit about what I'm now
claiming and what I'm withdrawing:

  • Withdrawn: "0.39x at 32 threads" (a harness artefact) and "yours is 9–11% faster at
    high thread counts" (my padding bug, now fixed).
  • Stands: at 100–400 visits and 16–64 threads, mine is 18–47% faster — measured before
    the fix, so if anything that understates it now.
  • New: at 3200 visits / 192 threads, mine is 1.23x yours with bucketing on at the
    shared default cap, or 1.33x comparing each of us at our own best nnMaxBatchSize.
    (Bucketing improves my own throughput 1.37x; that is a before/after on my backend, not a
    margin over yours.)

Caveat I'd rather state than have found: the low-visits grid predates the bucketing change,
so those cells deserve a rerun before anyone leans on them, and everything here is one
machine (MI325X, ROCm 7.2.0) with one net.

Thank you for pushing on the low-visits case. Without it I'd have published a number that
was wrong in your favour, and I'd never have looked at the padding that was costing my own
backend 1.37x.

@zhihuidu-amd

Copy link
Copy Markdown
Author

@Looong01 Two follow-ups: the low-visits grid you asked for, re-measured now that bucketing
exists, and one correction to my last comment.

The low-visits grid, redone — and the old one was wrong in my favour

Last time I flagged that the 100–400 visit numbers predated the bucketing fix and deserved a
rerun. They did, and the rerun changed more than I expected.

This time I ran three arms interleaved in one job on one node, each thread count in its
own process, warm-up discarded per cell:

  • #1234 — your ROCm backend, unchanged
  • ours-offmigraphxBatchBuckets=false, i.e. exactly the configuration I published last
    time
  • ours-on — bucketing enabled (my current default)

The ours-off arm is the point of the exercise. It re-measures the published configuration on
the same node in the same job, so any difference in ours-on is bucketing rather than drift.

visits threads #1234 ours-off ours-on off/1234 on/1234
100 16 1409.8 1870.5 1898.9 1.33x 1.35x
100 32 1481.0 2569.1 3012.0 1.73x 2.03x
100 64 2152.3 2874.1 4342.9 1.34x 2.02x
400 16 1778.2 2082.8 2111.3 1.17x 1.19x
400 32 3061.2 3041.9 3580.9 0.99x 1.17x
400 64 3209.3 3960.3 5220.9 1.23x 1.63x
1600 16 2049.7 2143.2 2183.8 1.05x 1.07x
1600 32 3697.6 3192.5 3717.6 0.86x 1.01x
1600 64 4526.6 4360.5 5497.4 0.96x 1.21x
3200 16 2129.3 2165.1 2192.1 1.02x 1.03x
3200 32 3718.7 3248.7 3745.6 0.87x 1.01x
3200 64 4686.4 4468.3 5564.2 0.95x 1.19x

avgBatchSize agrees across all three arms to within 0.5% in every cell, so the arms are doing
the same work.

Please treat this table as replacing the earlier one, not extending it. Your backend
measured materially faster in this job than in the original run — 1481 vs 1291 at v=100 t=32,
3061 vs 2554 at v=400 t=32 — so some of my published ratios were flattering to me. I don't
have a clean explanation for the shift, which is itself a reason not to mix the two sets. Had
I run only ours-on and compared it against the old #1234 column, every ratio in this table
would have been inflated.

Bucketing's benefit tracks thread count, not visits — about 1.5% at t=16, 16% at t=32,
24–51% at t=64. That is the padding mechanism showing itself: more threads means a larger
compiled shape relative to the batch the search actually produces.

So the answer to what you actually asked is yes, static-shape padding was hurting these
workloads
— every cell where I was behind (0.86x, 0.87x, 0.95x, 0.96x) sits in the
ours-off column and is gone in ours-on. I had the axis wrong, though: I expected the
damage to show up at low visits, and it is really a function of thread count.

Caveat: single trial per cell. The direction is consistent across all 12 and the bucketing
trend is monotonic in thread count, but individual figures carry roughly the 0.5% run-to-run
sd I measured earlier, plus whatever produced the #1234 shift above. One machine (MI325X),
one net, ROCm 7.2.0.

Correction: the im2col hypothesis is dead, and I said so too weakly

In my last comment I wrote that the im2col suspect "doesn't hold". I want to state the
conclusion more plainly, because you called it plausible and it's worth closing properly
rather than leaving as an open lead:

  • No supported MIGraphX knob moves the Im2d2Col_v2 share off 9.3% — MIGRAPHX_ENABLE_NHWC,
    MIGRAPHX_ENABLE_WINOGRAD, and MIGRAPHX_DISABLE_MIOPEN_FUSION all leave it where it was.
  • NHWC is 2% slower, not faster.
  • ENABLE_WINOGRAD is gfx12-only as documented, and is simply inert on gfx942.

So the hypothesis is disproved rather than merely unsupported. What I can show is that the
knobs don't change the kernel choice; what I cannot show is that a different choice would be
faster, which is why I filed nothing with AMD on this one.

Separately, while chasing the padding problem I did find a genuine MIGraphX bug — Softplus
and Softsign both call shape::lens() unguarded while parsing, so any ONNX model containing
either op fails to parse once an input dimension has min != max. That has a fix and tests
and is going upstream to ROCm/AMDMIGraphX. It's the first thing standing between this backend
and a genuine dynamic-batch shape; whether anything else stands behind it I don't yet know,
since I've only cleared the parse. Bucketing is what I have in the meantime, and none of this
affects #1234.

@Looong01

Copy link
Copy Markdown
Contributor

@Looong01 Two follow-ups: the low-visits grid you asked for, re-measured now that bucketing exists, and one correction to my last comment.

The low-visits grid, redone — and the old one was wrong in my favour

Last time I flagged that the 100–400 visit numbers predated the bucketing fix and deserved a rerun. They did, and the rerun changed more than I expected.

This time I ran three arms interleaved in one job on one node, each thread count in its own process, warm-up discarded per cell:

  • #1234 — your ROCm backend, unchanged
  • ours-off — , i.e. exactly the configuration I published last
    time
    migraphxBatchBuckets=false
  • ours-on — bucketing enabled (my current default)

The arm is the point of the exercise. It re-measures the published configuration on the same node in the same job, so any difference in is bucketing rather than drift.ours-off``ours-on

visits threads #1234 ours-off ours-on off/1234 on/1234
100 16 1409.8 1870.5 1898.9 1.33x 1.35x
100 32 1481.0 2569.1 3012.0 1.73x 2.03x
100 64 2152.3 2874.1 4342.9 1.34x 2.02x
400 16 1778.2 2082.8 2111.3 1.17x 1.19x
400 32 3061.2 3041.9 3580.9 0.99x 1.17x
400 64 3209.3 3960.3 5220.9 1.23x 1.63x
1600 16 2049.7 2143.2 2183.8 1.05x 1.07x
1600 32 3697.6 3192.5 3717.6 0.86x 1.01x
1600 64 4526.6 4360.5 5497.4 0.96x 1.21x
3200 16 2129.3 2165.1 2192.1 1.02x 1.03x
3200 32 3718.7 3248.7 3745.6 0.87x 1.01x
3200 64 4686.4 4468.3 5564.2 0.95x 1.19x
avgBatchSize agrees across all three arms to within 0.5% in every cell, so the arms are doing the same work.

Please treat this table as replacing the earlier one, not extending it. Your backend measured materially faster in this job than in the original run — 1481 vs 1291 at v=100 t=32, 3061 vs 2554 at v=400 t=32 — so some of my published ratios were flattering to me. I don't have a clean explanation for the shift, which is itself a reason not to mix the two sets. Had I run only and compared it against the old column, every ratio in this table would have been inflated.ours-on``#1234

Bucketing's benefit tracks thread count, not visits — about 1.5% at t=16, 16% at t=32, 24–51% at t=64. That is the padding mechanism showing itself: more threads means a larger compiled shape relative to the batch the search actually produces.

So the answer to what you actually asked is yes, static-shape padding was hurting these workloads — every cell where I was behind (0.86x, 0.87x, 0.95x, 0.96x) sits in the column and is gone in . I had the axis wrong, though: I expected the damage to show up at low visits, and it is really a function of thread count.ours-off``ours-on

Caveat: single trial per cell. The direction is consistent across all 12 and the bucketing trend is monotonic in thread count, but individual figures carry roughly the 0.5% run-to-run sd I measured earlier, plus whatever produced the shift above. One machine (MI325X), one net, ROCm 7.2.0.#1234

Correction: the im2col hypothesis is dead, and I said so too weakly

In my last comment I wrote that the im2col suspect "doesn't hold". I want to state the conclusion more plainly, because you called it plausible and it's worth closing properly rather than leaving as an open lead:

  • No supported MIGraphX knob moves the share off 9.3% — ,
    , and all leave it where it was.Im2d2Col_v2``MIGRAPHX_ENABLE_NHWC``MIGRAPHX_ENABLE_WINOGRAD``MIGRAPHX_DISABLE_MIOPEN_FUSION
  • NHWC is 2% slower, not faster.
  • ENABLE_WINOGRAD is gfx12-only as documented, and is simply inert on gfx942.

So the hypothesis is disproved rather than merely unsupported. What I can show is that the knobs don't change the kernel choice; what I cannot show is that a different choice would be faster, which is why I filed nothing with AMD on this one.

Separately, while chasing the padding problem I did find a genuine MIGraphX bug — and both call unguarded while parsing, so any ONNX model containing either op fails to parse once an input dimension has . That has a fix and tests and is going upstream to ROCm/AMDMIGraphX. It's the first thing standing between this backend and a genuine dynamic-batch shape; whether anything else stands behind it I don't yet know, since I've only cleared the parse. Bucketing is what I have in the meantime, and none of this affects #1234.Softplus``Softsign``shape::lens()``min != max

@zhihuidu-amd Thanks for running this properly - the three-arm interleaved design with the ours-off control arm is exactly right, and I accept the new table as replacing the old one. For the record: the drift you flagged (1481 vs 1291 etc.) is real and worth its own caveat on any cross-job comparison; your handling of it is more honest than most benchmark reports I've seen.

I'll also plainly retract what I said in #1188 about dropping MIGraphX. The bucketing data is convincing: with bucketing on, you win or tie every cell, and the 2x at 100v/64t is not noise-sized. My earlier "padding hurts at low visits" prediction was right about the mechanism existing but wrong about the axis (thread count, not visits) and wrong about the conclusion once bucketing exists. And on the im2col question you were right to close it - I called it plausible, your knob sweep disproved it. Noted and accepted.

The coherent picture as I now read it: MIOpen's per-op path wins when GPU-bound (big batches, high threads - your first 160/192-thread finding), the graph-compiler path wins when dispatch-bound (small batches, low-to-mid threads). That's exactly the CUDA vs TensorRT split on the NVIDIA side, and I no longer object to carrying both.

One important scope note, though: everything we've both measured so far is on b18c384nbt, a pure convnet. That model form is legacy as of v1.17 - the main training run is switching to transformers, and the three new strong nets (b10c384h6nbttflrs, b10c512h8nbt3tflrs-fson-silu-rsnh, b11c768h12nbt3tflrs-fson-silu) are all transformers. The transformer comparison between our two backends is still a blank page, and right now it isn't a level comparison:

  • On the ROCm side, transformer attention gets the CK FMHA fused-attention path on CDNA and RDNA3/3.5/4 (measured ~2x on nnEvals/s for the attention blocks on gfx1100 vs the builtin kernel), and the whole transformer trunk is NHWC-native with hand-written kernels.
  • On the MIGraphX side, your own finding stands unresolved: transformer nets under NHWC produce badly wrong policy output (policySqErr 136 vs 6e-10). Until that's root-caused, MIGraphX has to run transformers with the NCHW workaround, so any transformer benchmark would be measuring your safe fallback, not your best path.

So if you're willing, the most informative next measurement would be one of the new transformer nets (say b10c384h6nbttflrs for speed, or b11c768h12nbt3tflrs-fson-silu for the flagship) - same interleaved setup, plus a correctness spot-check of policy output against the Eigen or CUDA reference, since on transformers a backend can be fast and quietly wrong (we just caught exactly that class of bug on our own older ROCm branch the hard way). If the NHWC policy bug turns out to be in MIGraphX's lowering, it may be worth filing alongside your Softplus/Softsign find.

And thanks again for the rigor here - this is how backend comparisons should be done.

@yaoliu13

Copy link
Copy Markdown

@Looong01 Good to see you again in another repo. Thank you for your KataGo work. Really impressive!

Supporting multiple backends allows us to keep improving performance like #1237

Users can pick the one that fits their workload and GPU. What do you think?

@Looong01

Copy link
Copy Markdown
Contributor

@Looong01 Good to see you again in another repo. Thank you for your KataGo work. Really impressive!

Supporting multiple backends allows us to keep improving performance like #1237

Users can pick the one that fits their workload and GPU. What do you think?

Hii, thank u and good to see u. Do u have any suggestions in detail? Like what backends?

@zhihuidu-amd

Copy link
Copy Markdown
Author

@Looong01 Ran the transformer comparison you asked for. Correctness first, since your warning
about a backend being "fast and quietly wrong" is the more important half.

MI325X (gfx942), ROCm 7.2.0, b10c384h6nbttflrs from the v1.17.1 release — the "for speed"
net you suggested. Both backends built and run in one job on one node.

Caveat up front, because it shapes everything below: my arm is on the NCHW fallback.
migraphxTransformerNHWC still defaults false, so this is not my best path — exactly as you
pointed out. You're comparing your NHWC-native hand-written kernels with CK FMHA against my
safe fallback.

Correctness

runnnonmanyposestest over the bundled position set, 92202 output values per arm
(win probs, score means, policy probs). Reference is my FP32 output:

arm winProbSquerr scoreMeanSquerr policyProbSquerr
mine, FP16 1.37e-03 2.98e-01 5.08e-03
#1234, FP32 4.18e-11 2.96e-08 2.42e-10
#1234, FP16 5.94e-05 2.01e-02 4.65e-04

Two independent implementations — a graph compiled by MIGraphX versus your hand-written
HIP/MIOpen kernels — agreeing to ~1e-10 across 92202 values is the result I'd hoped for.
Nothing resembling the policySqErr 136 that the NHWC path produces.

One result goes against me and I'd rather flag it than have you find it: in FP16 my error
is ~23x yours on winProb and ~21x on policy. On the convnet my FP16 was more accurate than
OpenCL's, so this is the opposite direction. Both are far below anything that could change a
move choice — 5e-3 summed squared error over 92202 policy values — but I don't have an
explanation yet and I'm not claiming FP16 parity on transformers until I do.

Comparing both GPU backends against my FP32 only proves we agree, not that we're both right,
so I built the Eigen CPU backend and re-ran against that — an independent implementation, and
the check you actually asked for:

arm winProbSquerr scoreMeanSquerr policyProbSquerr
mine, FP32 vs Eigen 1.17e-09 9.77e-07 1.62e-09
#1234, FP32 vs Eigen 1.31e-09 1.18e-06 1.67e-09
mine, FP16 vs Eigen 1.44e-03 3.01e-01 4.90e-03
#1234, FP16 vs Eigen 5.93e-05 2.01e-02 4.65e-04

Both FP32 paths land within ~1e-9 of a CPU implementation that shares no code with either.
That closes the "fast and quietly wrong" question for FP32 on this net. (All arms run
inputsUseNHWC=true here — the Eigen backend rejects false, and mixing layouts across arms
would compare different tensors.)

The FP16 gap survives the independent reference unchanged, so it is real and mine, not an
artefact of using my own output as the baseline.

Throughput

Same protocol as the low-visits grid: each thread count in its own process, warm-up discarded
per cell, arms interleaved. Single trial per cell.

visits threads #1234 mine (NCHW) ratio
400 16 1071.2 2474.7 2.31x
400 32 1650.8 3839.3 2.33x
400 64 2290.4 5100.3 2.23x
400 128 2956.3 5912.4 2.00x
400 192 3483.3 6539.4 1.88x
3200 16 1177.8 2567.6 2.18x
3200 32 1940.7 3997.8 2.06x
3200 64 2937.1 5074.4 1.73x
3200 128 4264.4 5841.6 1.37x
3200 192 4650.8 6017.1 1.29x

avgBatchSize matches within ~1% in 8 of 10 cells (worst 7%, at 3200/128), so the arms are
doing the same work.

What I read from it

Your dispatch-bound/GPU-bound synthesis holds, and the transformer shifts the whole curve.
The margin decays monotonically with batch size — 2.3x at bs≈8-15 down to 1.29x at bs≈93-100 —
the same shape as the convnet. But the level is much higher: at t=16 the convnet gap was
1.03-1.07x, here it's 2.18-2.31x.

I'll say the part that's awkward for my own framing: I expected your fused-attention path to
show up as an advantage somewhere in this table, and it doesn't. I'd guess CK FMHA is a real
win per attention block and is being swamped end-to-end by per-op dispatch at these batch
sizes — but that's a hypothesis, not something I measured, and you're better placed than I am
to say whether it's right.

The flagship net, where it goes the other way

I also ran b11c768h12nbt3tflrs-fson-silu. I expected its larger per-layer GEMMs to push
toward the GPU-bound regime where you do better, and they do:

visits threads #1234 mine (NCHW) ratio
400 16 589.5 989.9 1.68x
400 32 878.0 1362.2 1.55x
400 64 1012.5 1569.9 1.55x
400 128 1372.0 1725.5 1.26x
400 192 1518.8 1819.7 1.20x
3200 16 600.5 1014.1 1.69x
3200 32 953.2 1388.6 1.46x
3200 64 1327.6 1524.3 1.15x
3200 128 1607.7 1560.3 0.97x
3200 192 1727.8 1719.4 1.00x

Side by side at the corners: 2.31x → 1.68x at 400/16, and 1.29x → 0.97x at 3200/128.
You take the top-right corner on the flagship — the first cells either transformer has
gone against me, and they sit exactly where your model says per-op kernels should win.

Two caveats, one in each direction. At 0.97x/1.00x the gap is inside the run-to-run noise
I measured earlier (0.22-0.68% my arm, 1.1-2.5% yours), so "parity at the top end" is the
honest reading rather than "you win". Against that, my avgBatchSize in those two cells is
larger than yours — 67.3 vs 62.5 at t=128, 108.3 vs 93.7 at t=192 — so I'm doing ~16% more
work per batch for the same throughput at t=192, and reading that as a clean tie flatters me.

Correctness on the flagship matches the small net: #1234 FP32 vs mine at 1.58e-10, and my
FP16 error again larger than yours (~6x winProb, ~17x policy).

Limits worth stating: single trial per cell; one machine, one ROCm version, one net per size
class.

On the NHWC bug: migraphx-driver verify --gpu passes on both emitted graphs, so MIGraphX
computes the channel-last graph correctly and the fault is in my emitter, not something to
file against MIGraphX. That one's mine to fix, and until it is, the numbers above are what I
actually ship rather than the best I could do.

@yaoliu13

Copy link
Copy Markdown

@Looong01 Prefer to merge this PR now and put follow-ups in smaller PRs rather than expanding this one: https://testing.googleblog.com/2024/07/in-praise-of-small-pull-requests.html

OK to merge?

@yaoliu13

Copy link
Copy Markdown

@lightvector Could you review this PR? Thank you.

@yaoliu13

Copy link
Copy Markdown

We can run KataGo on AMD GPUs in a few different ways, similar to how NVIDIA supports OpenCL, cuDNN, and TensorRT. Supporting multiple options keeps the door open while we keep improving performance. Users can pick the one that fits their workload and GPU version.

@lightvector

Copy link
Copy Markdown
Owner

I'm reluctant to merge yet another backend. KataGo already has seven backends. This is already too much maintenance burden. Do you have suggestions on how to deal with this? For example, does this backend 100% obsolete the rocm backend that we just merged, such that we can delete it? @Looong01 - curious for your thoughts as well, given that it was your rocm backend that was just merged, and earlier you had worked on a migraphx backend that I had reviewed that you had claimed was significantly worse which was why we went with rocm instead.

@yaoliu13

Copy link
Copy Markdown

@lightvector @Looong01 Fair question on maintenance — we would not frame this as “replace ROCm.” The benchmark thread shows MIGraphX and ROCm trade off. The two paths look complementary, similar to TensorRT vs CUDA.

At AMD, we support multiple paths on ROCm depending on customer requests.

We’d suggest merging this PR, keeping ROCm, and putting performance improvement follow-ups in small PRs with clear docs on when to use each recipe. Happy to help with validation and maintenance from AMD’s side.

@lightvector

Copy link
Copy Markdown
Owner

@yaoliu13 Thanks. Can you explain more about the difference between this backend and #1222 which was merged just recently? If this backend is using ONNX anyways, and #1222 has migraphx as a provider for onnx's runtime, what is it doing differently than this PR? Is this PR more efficient somehow and what is the reason for the difference?

@Looong01

Looong01 commented Aug 16, 2026

Copy link
Copy Markdown
Contributor

I built the MIGraphX backend branch from https://github.com/zhihuidu-amd/KataGo/tree/migraphx-backend locally and ran the same four release models through both the MIGraphX backend and the ROCm backend on my RX 7900 XTX (gfx1100 ROCm 7.14, MIOpen 3.5.2. The GTP protocol works on MIGraphX (first move generated correctly), and the benchmark numbers below are the measured end-to-end visits/s from katago benchmark.

Comparison table

Model Visits Threads ROCm visits/s MIGraphX visits/s ROCm / MIGraphX
b10c384h6nbttflrs 400 16 1807 823 2.20x
b10c384h6nbttflrs 400 32 2326 818 2.84x
b10c512h8nbt3tflrs-fson-silu-rsnh 400 16 1222 466 2.62x
b10c512h8nbt3tflrs-fson-silu-rsnh 400 32 1433 413 3.47x
b11c768h12nbt3tflrs-fson-silu 400 16 747 250 2.99x
b11c768h12nbt3tflrs-fson-silu 400 32 779 233 3.35x
kata1-zhizi-b40c768nbt-s11272M-d5935M 400 16 316 269 1.18x

What the numbers show

  • ROCm is faster on every model and configuration. The margin is smallest on the largest pure-convolution network (b40,1.18x) and largest on the transformer models at higher thread counts (up to ~3.5x).
  • Thread scaling is very different. The ROCm backend benefits clearly from more threads (e.g. b10 rises from 1807 2326 visits/s going from t=16 to t=32; b11 rises from 747 to 779). The MIGraphX backend is essentially flat or slightly slower when thread count increases (b10: 823 → 818; b: 250 → 233). This matches the "static-shape graph compiled per bucket" hypothesis: once the graph is compiled for a fixed batch shape, throwing more CPU threads at the search does not help the GPU side if the compiled shape does not line up with the actual search batch.
  • The gap grows as the model becomes smaller and more dispatch-bound. b40 is heavy enough that the GPU stays busy regardless of the backend, so the difference is only ~18%. smaller transformer models are more sensitive to per-op dispatch/host overhead, and there the ROCm backend's hand-written kernels and dynamic batching pull ahead by a large margin.

Important caveats about this MIGraphX build

This is not the possible MIGraphX configuration. My local ROCm 7.14 apt repository does not provide the rocmlir or composable_kernel packages, so I had to build MIGraphX from the rocm-7.14 with:

  • MIGRAPHX_ENABLE_MLIR=Off
  • MIGRAPHX_USE_COMPOSABLEKERNEL=Off
  • MIGRAPHX_USE_HIPBLASLT=Off (the installed hipBLASLt package is missing the gfx1100 Tensile library files)

I also had to apply small local patches to make MIGraphX compile at all in this configuration: stubs for is_module_fusible, adjust_param_shapes, dump_mlir_to_file, dump_mlir_to_mxr (which are only compiled when rocMLIR is enabled), and a fix for gfx_default_rocblas() being undefined when hipBLASLt is disabled.

So the MIGraphX arm here is effectively IOpen for convolutions + rocBLAS for GEMMs, no rocMLIR fusion, no CK. That is a lower-bound performance for MIGraphX; the published numbers from your setup may be noticeably better if you have rocMLIR working. Even with that caveat, the ROCm backend is still faster on the same hardware, and the thread-scaling pattern is striking.

Bottom line

The MIGraphX backend compiles and runs correctly, but my RDNA3 machine it does not match the ROCm backend's throughput or scaling. The difference is modest on the heaviest network (b40) and large on the smaller transformer networks where overhead dominates. I think the ROCm backend should remain the primary AMD GPU path in KataGo, and MIGraphX is better kept as a separate experimental branch until it can match this performance or provide a clear feature advantage.

@Looong01

Looong01 commented Aug 16, 2026

Copy link
Copy Markdown
Contributor

Follow-up: testing the "dispatch-bound" hypothesis with hipGraph capture

I want to refine one point from my earlier numbers, because your transformer comment included a specific hypothesis:

"I'd guess CK FMHA is a real win per attention block and is being swamped end-to-end by per-op dispatch at these batch sizes."

I ran a direct experiment to test this: capture the whole neural-net forward pass into a single HIP graph and replay it, collapsing the 100+ individual HIP dispatches down to one graph launch. This is the cleanest way to isolate dispatch overhead from kernel execution time.

On gfx1100 (RX 7900 XTX) the result is neutral within ~3% run-to-run noise. In other words, on this RDNA3 card, per-op dispatch overhead is not the bottleneck. If CK FMHA were a real per-block win but being hidden by dispatch, hipGraph would have exposed it; it did not.

So the MIGraphX 2× advantage on gfx942 (MI325X) is probably not coming from "fewer dispatches" in the graph-compiler sense. A sharper explanation is that MIGraphX's op fusion reduces the number of kernels and the amount of device-memory traffic, and at small batch sizes those kernels are tiny and memory-latency-dominated. That is a different bottleneck than dispatch count.

Architecture asymmetry caveat

This experiment only holds for RDNA3. CDNA (gfx942) has a different driver/firmware dispatch path, and dispatch overhead may genuinely be more expensive there. I do not have a CDNA machine, so I cannot run the same hipGraph check on MI325X. Someone with CDNA access would need to repeat the capture test to settle whether the two architectures have different bottlenecks.

Scorecard summary after the transformer run

  • Correctness: both backends agree with Eigen CPU at ~1e-9 in FP32. Closed.
  • FP16 accuracy: our ROCm backend's error is ~20× smaller than yours, which is a real advantage of FP32-accumulation GEMM.
  • NHWC bug: you confirmed it is in your emitter, not in MIGraphX itself. Once fixed, your NHWC transformer numbers will improve.
  • Throughput: the advantage regions still hold. Small-batch transformer workloads favor MIGraphX by a large margin; large-batch flagship workloads are at parity or favor ROCm once effective work is accounted for.

My reading is unchanged: the two backends win in different regions, and neither is universally better. The ROCm backend should stay as the default AMD path because it is already faster on the same hardware in most KataGo use cases and matches or exceeds MIGraphX on the largest models at high batch sizes.

@Looong01

Looong01 commented Aug 16, 2026

Copy link
Copy Markdown
Contributor

On my side, as my result:

Maybe MIGraphX backend runs faster than ROCm backend with CK on CDNA machine, but I only have RDNA machine so I can NOT prove @zhihuidu-amd's result on my side. Maybe anyone who have CDNA resources can help. In fact, ROCm runs better on my RDNA machine than MIGraphX.

@yaoliu13

Copy link
Copy Markdown

@lightvector #1222 mentioned that other providers (cuda / migraphx / coreml) are wired in and should work but are unverified, and the code stated that MIGraphX from-source build slots in here once validated (needs ROCm; deferred). Some users may run ONNX with MIGraphX but some may prefer running MIGraphX directly.

@Looong01 MIGraphX officially supports MI300X on Ubuntu [1] but doesn't officially support RX 7900 XTX (gfx1100) on Windows [2].

1: https://rocm.docs.amd.com/en/latest/compatibility/compatibility-matrix.html?fam=instinct&w=compute&gpu=mi300x&gfx=gfx942&os=ubuntu
2: https://rocm.docs.amd.com/en/latest/compatibility/compatibility-matrix.html?fam=radeon&w=compute&gpu=rx-7900-xtx&gfx=gfx1100&os=windows

@lightvector

Copy link
Copy Markdown
Owner

Thanks, so if migraphx is already available under the onnx backend that was merged, then my request would be to test and make whatever adjustments are necessary to get migraphx working well there instead of merging this PR, which seems to also a wrapper around onnx. Is that possible or am I missing something?

Dealing with all the all the different accelerator providers each requiring custom implementations is a big burden, and I'd like to be moving towards unifying under a smaller number of common frameworks. If TensorRT weren't already grandfathered in due to the historical development path, I'd also be pushing for TensorRT to be available via only the onnx backend if possible rather than being its own backend.

@yaoliu13

Copy link
Copy Markdown

@lightvector This PR does not use ONNX Runtime, and it does not load a pre-built .onnx file. We use ONNX only as an in-memory intermediate graph format — built by onnxmodelbuilder.cpp / OnnxModelBuilder::build(), the same emitter the TensorRT backend uses — and then pass that buffer directly to MIGraphX via parse_onnx_buffer. That mirrors how the TensorRT backend works today (ONNX emitter → TensorRT via nvonnxparser), without going through ONNX Runtime.

We’re happy to investigate getting MIGraphX working under the ONNX backend (#1222) as you suggested. For context, https://rocm.blogs.amd.com/artificial-intelligence/triton-inference-server/README.html benchmarks ONNX with MIGraphX vs. ONNX with TensorRT. For fair comparisons, we should align on the serving stack. Otherwise, comparing peak-performance native TensorRT on one side against ONNX + MIGraphX (not native) on the other wouldn’t be apples-to-apples. We should also allow users to use native MIGraphX for flexibility.

@zhihuidu-amd

Copy link
Copy Markdown
Author

@lightvector On the #1222 question — I built the comparison arm rather than argue from the code,
and the result changes the picture, so I want to lead with it.

ONNX Runtime + migraphx EP, measured

ORT built from source at v1.20.1 with --use_migraphx; KataGo built USE_BACKEND=ONNX from the
same merged tree, same node, as the other two arms. MI325X, ROCm 7.2.0, b18c384.

arm provider threads visits/s
onnx backend migraphx 1 273.67
onnx backend migraphx 16 0.36
onnx backend cpu 1 0.72

The EP works — at threads=1 it is ~380x the CPU provider, so the graph really is running on the
GPU. But going from 1 thread to 16 costs 760x.

The cause is in ORT's own source
(onnxruntime/core/providers/migraphx/migraphx_execution_provider.cc, v1.20.1 line 1298):

// input shapes are different, needs to re-parse onnx and
// re-compile the program
if (!input_shape_match) {

MIGraphX compiles static shapes. MCTS submits a different batch size on nearly every eval — that
is the point of the search — so the EP re-parses the ONNX and re-runs the MIGraphX compiler
inside the inference call
, continuously. Compilation is seconds; inference is milliseconds. At
threads=1 the batch is pinned to 1, one shape is ever seen, the graph compiles once, and it is
fast.

Is there a config fix? Not in the version a #1222 user would build today. ORT 1.27 added a
shape-keyed .mxr cache (ORT_MIGRAPHX_MODEL_CACHE_PATH). v1.20.1 has nothing equivalent —
grep -c model_cache on that file returns 0; it only has save/load against a single fixed path,
which cannot serve a workload with many shapes. I set the 1.27 variable against the 1.20.1 build
anyway and the cache directory stayed empty, consistent with the knob not existing there.

(I nearly got this wrong in your favour: I first read the recompile code in a 1.27 checkout that
happened to be on our cluster and almost attributed 1.27's caching to the build I had actually
measured. The versions differ in exactly the way that matters.)

What "get migraphx working well under #1222" would actually take

Your instinct to unify is right, and I am not going to argue against it on weak grounds. But it is
not a config change or a docs line. It needs one of:

  1. ORT >= 1.27, plus wiring ORT_MIGRAPHX_MODEL_CACHE_PATH from a KataGo config key, plus
    confirming the shape-keyed cache actually amortizes under MCTS's shape distribution. I have not
    measured that and will not assert it.
  2. Batch-shape stabilization inside the ONNX backend — padding every eval up to a small ladder
    of fixed shapes so the EP sees few distinct shapes.
  3. Upstream work in the MIGraphX EP so it keeps a per-shape program cache in memory.

Option 2 is worth dwelling on, because it is the same problem this PR already solves. The reason
this backend compiles a geometric ladder of batch buckets is precisely that MIGraphX compiles
static shapes. #1222 hits the identical wall and currently handles it by recompiling. Any fix
under #1222 converges on the bucketing logic that is already written and tested in this PR.

So the two paths are not really "new backend vs. reuse the ONNX one" — they are "this logic lives
in a MIGraphX backend" or "this logic gets ported into the ONNX backend." I am happy to do the
second if that is what you prefer; it is the same engineering either way, and I would rather do the
version you will merge. Tell me which and I will build it.

Two other differences worth knowing, since both are ONNX-backend gaps rather than design limits:


@Looong01 Thank you for actually building and running it — that is more than I had any right to
expect, and the RDNA3 data point is one I cannot produce.

One thing in your setup needs flagging, and it is what your own caveat already names. You built
MIGraphX with MIGRAPHX_ENABLE_MLIR=Off, MIGRAPHX_USE_COMPOSABLEKERNEL=Off, and
MIGRAPHX_USE_HIPBLASLT=Off. Those are where essentially all of MIGraphX's fusion lives. Here is
what my build actually does on this net (MIGRAPHX_TRACE_MLIR=1, b10c384h6nbttflrs, bs=32):

    339  mlir_convolution
     85  mlir_transpose_slice_reshape_transpose_reshape_dot_mul_convert_reshape_
         reduce_max_reshape_sub_exp_reshape_reduce_sum_reshape_div_convert_dot
     69  mlir_convolution_convert
     63  mlir_convolution_add
     40  mlir_convolution_mul_add_add_sigmoid_mul

The second line is an entire attention block — Q·Kᵀ, scale, softmax, ·V — fused into one
kernel. The last is conv+bias+SiLU. With MLIR off, each becomes separate MIOpen/rocBLAS calls with
full device-memory round trips between them.

Rather than assert that matters, I measured it — same machine, same graphs, isolated compiled
program, toggling only MIGRAPHX_DISABLE_MLIR:

net batch MLIR on MLIR off cost
b10c384h6nbttflrs 32 501.6 inf/s 373.2 1.34x
b10c384h6nbttflrs 128 453.5 349.6 1.30x
b11c768h12nbt3tflrs-fson-silu 32 264.8 169.6 1.56x
b11c768h12nbt3tflrs-fson-silu 128 341.8 241.5 1.42x

That is MLIR alone on gfx942; your build also had CK and hipBLASLt off, so the real handicap is
larger. Note the shape of it: the penalty is worst on the transformers (1.56x), which is exactly
where your table shows your largest margins (up to 3.5x), and mildest on the convnet, matching your
smallest (1.18x on b40). That correlation is why I think a good part of what the table measures is
the missing fusion stack.

Your writeup already calls this "a lower-bound performance for MIGraphX" — I would just ask that the
caveat travel with the table wherever it is cited, because at a glance it reads as a like-for-like
comparison.

To be explicit about what I am not claiming: I am not saying MIGraphX would win on RDNA3 with a
full build. I have no RDNA3 hardware — ours is CDNA3 (MI300X/MI325X, gfx942), a different
architecture with different MLIR maturity. It may well still lose on gfx1100. I am only saying this
particular comparison cannot settle it either way.

My own numbers, including the ones against me

Head-to-head vs merged master ROCM (not the #1234 branch), one tree, one node, one job, 5
interleaved trials per cell, sd < 2% everywhere:

net threads this PR rocm ratio
b10c384h6nbttflrs 16 2397.5 1169.5 2.05x
b10c384h6nbttflrs 32 3276.4 1942.9 1.69x
b10c384h6nbttflrs 64 3966.0 3082.9 1.29x
b10c384h6nbttflrs 192 4654.9 4452.9 1.05x
b11c768h12nbt3tflrs-fson-silu 32 995.9 929.1 1.07x
b11c768h12nbt3tflrs-fson-silu 64 1097.5 1314.5 0.83x
b11c768h12nbt3tflrs-fson-silu 128 1150.9 1576.0 0.73x
b11c768h12nbt3tflrs-fson-silu 192 1235.1 1683.0 0.73x
b18c384 128 5387.7 5587.1 0.96x
b18c384 192 5668.5 6106.2 0.93x

Two corrections to my own published numbers, both against me:

  1. The 1.22x I posted earlier is stale. Measured against the Backends to merge #1234 branch; against merged master
    the same cell is 0.92x. You shipped a lot of fixes after the merge and they landed.
  2. Your FP16 accuracy advantage holds against an independent reference. Versus Eigen CPU over
    92202 values: policy squerr 2.12e-04 (yours) vs 4.10e-03 (mine) on b18c384; 4.65e-04 vs 5.25e-03
    on the transformer. 6.5x and 22x, your way. Unexplained on my side, and I am not claiming FP16
    parity until it is.

@lightvector So the direct answer to your original question is no — this does not 100% obsolete the
ROCm backend, and I would not ask you to delete it on the strength of these numbers. We lead below
batch ~50 and trail above it.

What I would ask instead: the small-batch region is not a corner case, it is analysis, GTP play, and
anything latency-sensitive, and 2.05x there is worth having. If the blocker is backend count rather
than the measurements, I would rather adapt this PR than drop the work — porting the bucketing into
the ONNX backend, wiring the EP's FP16 flag, or whatever shape makes it mergeable for you. Tell me
which direction is acceptable and I will do that work.

@Looong01

Looong01 commented Aug 17, 2026

Copy link
Copy Markdown
Contributor

@lightvector #1222 mentioned that other providers (cuda / migraphx / coreml) are wired in and should work but are unverified, and the code stated that MIGraphX from-source build slots in here once validated (needs ROCm; deferred). Some users may run ONNX with MIGraphX but some may prefer running MIGraphX directly.

@Looong01 MIGraphX officially supports MI300X on Ubuntu [1] but doesn't officially support RX 7900 XTX (gfx1100) on Windows [2].

1: https://rocm.docs.amd.com/en/latest/compatibility/compatibility-matrix.html?fam=instinct&w=compute&gpu=mi300x&gfx=gfx942&os=ubuntu 2: https://rocm.docs.amd.com/en/latest/compatibility/compatibility-matrix.html?fam=radeon&w=compute&gpu=rx-7900-xtx&gfx=gfx1100&os=windows

All my tests are on Ubuntu, not Windows. Btw, there is another reason to support ROCm, that ROCm support both Linux and Windows, but MIGraphX only support Linux.

@Looong01

Looong01 commented Aug 17, 2026

Copy link
Copy Markdown
Contributor

I have no ill intentions towards adding a new backend; I just want to discuss the issue itself. @lightvector mentioned that his current maintenance workload is quite heavy.

I want to step back and ask what user population a merged MIGraphX backend would actually serve.

The use case is CDNA training, not RDNA inference

Your small-batch wins (b10/t16 2.05×, b10/t32 1.69×, b11/t32 1.07×) are measured on MI300X/MI325X (gfx942). At large batch and on larger models the same table shows ROCm winning. That means the MIGraphX advantage region is specifically:

  • CDNA / Instinct hardware
  • relatively small batch sizes
  • workloads that care about throughput at small batch

In KataGo, the only workload that fits all three is large-scale selfplay / training data generation. Individual users running GTP or analysis on a single workstation do not drive enough small-batch evals to make that 2× region decisive; they are more likely to be memory- or large-batch-limited, where the ROCm backend is already faster or equal.

But training is not on the table

The public training pipeline at https://katagotraining.org/ is CUDA-based. As far as I know, @lightvector has not announced plans to move KataGo training to AMD GPUs, and neither ROCm nor MIGraphX backends are part of that infrastructure today. If the MIGraphX backend is not going to be used for KataGo's own distributed training, then its natural use case disappears.

Most end users are on RDNA

The people who actually download KataGo binaries and run them on AMD hardware are overwhelmingly on consumer RDNA cards (RX 6800/6900/7800/7900 series). On that hardware:

  • ROCm is officially supported on both Linux and Windows.
  • MIGraphX is not officially supported on Windows for RX 7900 XTX, and the Linux packages for full fusion (migraphx) are not in the standard ROCm 7.14 consumer repo(Just 7.2 has it).
  • Our measurements on gfx1100 show ROCm ahead across the board, with the MIGraphX build handicapped by missing MLIR/CK/hipBLASLt.

So for the user base that actually exists today, the ROCm backend is the right path. A merged MIGraphX backend would primarily serve a hypothetical future where KataGo training runs on CDNA clusters, and that future is not currently planned.

If, after considering all the factors I've mentioned above, @lightvector still believes the MIGraphX backend deserves to be merged, then I have no further comments. I fully respect our repo owner's ideas, plans, and arrangements.

@lightvector

lightvector commented Aug 17, 2026

Copy link
Copy Markdown
Owner

@zhihuidu-amd thanks for the new report and benchmarks! These are extremely helpful to understand the direction... which I think does actually suggest not merging this backend.

@Looong01 thanks for the input and nice argument, but if your post is AI-written, also consider extra-double-checking the AI-written posts or try to have them be less overconfident about places they might be making claims that are inaccurate.

  • CDNA / Instinct hardware
  • relatively small batch sizes
  • workloads that care about throughput at small batch

In KataGo, the only workload that fits all three is large-scale selfplay / training data generation.

Selfplay data generation is large-batch-size, not small batch size, to benefit from the greater throughput of large batch by running a large number of games in parallel.

The public training pipeline at https://katagotraining.org/ is CUDA-based.

Not quite. The data generation pipeline (which is the vast majority of the compute) is whatever hardware all the contributors running KataGo are using to generate and upload data. There is nothing that ties it to CUDA, although it is true that in practice NVIDIA gpus are more common among users at the moment. The minority of the compute, the actual training of the neural net itself, is pytorch, which is also technically cuda in practice on the particular GPU machines we use for it but not relevant to a discussion of the C++ backend.

Those are the two major inaccuracies I noticed.

But in any case, @Looong01 correcting the first inaccuracy above actually improves your argument. Since even self-play data generation does not fit the above advantage profile (being large-batch), and if indeed RDNA is what the vast majority of consumers have, that would leave the advantage region of MIGraphX without any use case.

I'll note that also b11c768h12nbt3tflrs-fson-silu is unambiguously the best model even considering compute cost, and there the benchmark showed rocm was better. The only reason for releasing the small transformers (on which MIGraphX appeared to be better) was to offer an alternative for very weak hardware, but the small transformers are unambiguously weaker per compute cost, the only reason you would run them is if your hardware was so weak that you simply could not run the larger models to a minimal number of visits at all, or for research purposes into suboptimally small models. If indeed CDNA is the "datacenter" generation of devices, then I would be surprised if users would be going out of their way to rent datacenter GPUs only to run a weak small model. Thoughts? Is there some case where MIGraphX would still add a lot of value?

@zhihuidu-amd

Copy link
Copy Markdown
Author

@lightvector One important finding before those benchmarks are used to draw a conclusion: we
found a configuration error in how this backend sizes its compiled graph, and with it fixed the
backend is faster than ROCm on the cells we have re-measured so far
— including the flagship net
you singled out.

The configuration error

MIGraphX compiles a static shape. This backend compiles one program at nnMaxBatchSize and
zero-pads every evaluation up to it. The design note in the file asserts "MCTS batches are
near-full in practice." I finally measured that assumption instead of trusting it:

threads compiled shape measured avgBatchSize wasted compute
16 16 7.96 2.01x
32 32 15.82 2.02x
64 64 31.47 2.03x
128 128 63.63 2.01x
192 192 101.02 1.90x

MCTS fills almost exactly half the batch, at every thread count. So the benchmarks I posted had
this backend computing ~2x the rows it needed, while the ROCm backend — which uses dynamic shapes —
paid nothing for this. It also explains something I had reported as unexplained: the isolated
compiled program benchmarked ~1.9x faster than the same program running inside the application.
Same factor, same cause.

The fix is batch bucketing: compile a small geometric ladder of shapes and dispatch each eval to
the smallest that fits. It had been written and validated earlier but was not present on this
branch, so every benchmark posted so far was built without it. Now restored and re-measured.

The results I have so far

Not projections — measured. But partial, and I am labelling exactly how partial.

1. Bucketing on vs off vs ROCm, on merged master (b11c768h12nbt3tflrs-fson-silu, t=64,
visits=3200, 3 interleaved trials, same node/tree/job):

arm nnEvals/s vs ROCm
this PR, bucketing on 1509.5 1.15x
this PR, bucketing off (the earlier configuration) 1110.9 0.84x
rocm (unchanged binary) 1315.8

A 0.83x loss becomes a 1.15x win; bucketing alone is worth 1.36x to this backend.

2. Bucketing on vs off at larger batch (b18c384, visits=3200, 5 interleaved trials, our
backend both arms — this is the older tree, before the merge, so treat it as the size of the
effect rather than as a head-to-head):

threads bucketing off bucketing on gain
96 4995.1 6550.9 1.31x
192 5444.8 7453.2 1.37x

The gain does not fade at high thread counts, which is the relevant question for your point about
large-batch workloads.

3. Correctness, checked before any timing — policy squerr 1.8e-10 against the
Eigen-verified FP32 reference. Bucketing is not buying speed with accuracy.

What I am not claiming yet

I do not know that this backend now wins everywhere, and I would rather say so than find out in
public. A full 5-trial re-run of all ten cells from my previous table is running now — same
protocol, same node, ROCm binary unchanged as the control. I will post the complete table,
including any cell where this backend still loses.
The open question is the two largest-thread
flagship cells: padding waste there was 1.90x rather than ~2.03x, so there is slightly less to
recover, and an earlier compiled-shape sweep showed throughput on that net peaking around bs=64
and declining above it.

Results in a few hours.

This bears directly on your remark that b11c768h12nbt3tflrs-fson-silu is the model that actually
matters. That is precisely the net in table 1 above — the one the earlier configuration showed ROCm
winning, and where this backend now leads by 1.15x.

Everything else in my previous comment stands unchanged — the ONNX-EP recompile finding, and
Looong01's FP16 accuracy advantage, which is real and still unexplained on my side.

@Looong01 Your point about Windows support on RDNA is fair and I cannot answer it with hardware I
do not have. One thing worth noting: the "MIGraphX only helps at small batch" conclusion that we
had both converged on appears to be an artifact of this bug rather than a property of the backend.
Whether that changes anything on gfx1100 I genuinely do not know — but the padding waste is a
function of MCTS batch-filling behaviour, not of the GPU, so it would have applied to your
measurements too.

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.

4 participants