HIP EP Enablement for MI350 by penglu12-quark · Pull Request #602 · ROCm/hip-ep

penglu12-quark · 2026-07-31T01:33:36Z

Summary

Enablement of hip-ep on MI350X (gfx950 / CDNA4 / wave64)

GQA softmax launches now take their block-size floor from the runtime wave size
instead of a hard-coded 64, which restores the original block geometry on
wave32.
The protobuf dllexport_decl workaround is narrowed to a new
MORPHIZEN_PROTO_DLL_SPEC, so it no longer strips the visibility attribute
from morphizen's hand-written public API headers.

Why

Softmax block floor. The wave64 port raised the softmax block-size floor to a
hard-coded 64 for a real reason: a 32-thread block is a partial wave on CDNA, and
the inactive upper lanes get folded into the __shfl_xor reduction. But the floor
is unconditional, so it also applies on RDNA, where it forces a second wave for
every row count <= 32 — the common decode shape — with nothing to show for it.
The floor needs to be the device's wave size, not the larger of the two.

num_waves has to be derived from that same value, because it sizes the dynamic
LDS fan-in buffer and must match what the kernel computes internally
(blockDim.x / WAVE_SIZE). Previously it used HIPDNN_WAVE_SIZE, which the host
pass hardcodes to 32, so on gfx950 that buffer was over-allocated by 2x.

Wave-size probe caching. The probe cached a single answer for the whole
process. A host can drive GPUs from different families at once (e.g. a gfx1151 APU
alongside a discrete card), where that cached answer gets applied to the wrong
device after hipSetDevice.

Export macro over-reach. protobuf v34 emits
class <decl> [[gnu::warn_unused]] Name, and GCC rejects a GNU __attribute__
placed before a C++11 [[...]] attribute, so the decoration has to be empty on
ELF. The previous fix achieved that by emptying MORPHIZEN_DLL_SPEC itself — but
that macro also decorates roughly thirty hand-written public API headers
(graph.hpp, pass.hpp and friends), not just generated code.

That is harmless today because nothing is built with -fvisibility=hidden, which
is exactly what makes it worth fixing now: it is a silent trap. If hidden
visibility ever reaches these targets, the public API stops being exported and the
only symptom is a link failure in a downstream consumer. Splitting the macro keeps
the workaround where it belongs and costs nothing.

What

hip_arch_compat.h: replace hipdnn_device_is_wave64() with
hipdnn_device_wave_size() returning 32 or 64, cached per device ordinal (up to
16) with an uncached probe beyond that. The cache array is zero-initialized so
no thread-safe-statics guard is emitted, which matters because this header is
compiled into bitcode with -fPIC. A failed query still yields the wave32
default. hipdnn_device_is_wave64() is kept as a thin wrapper.
gqa_kernel.hip: both softmax launchers (launch_softmax_f32_to_out,
hip_gqa_softmax_inplace) seed threads and compute num_waves from
hipdnn_device_wave_size().
morphizen/export.h: restore MORPHIZEN_DLL_SPEC to
__attribute__((visibility("default"))); add MORPHIZEN_PROTO_DLL_SPEC, empty
on ELF and __declspec(dllexport) on Windows.
morphizen-core/cmake/proto.cmake, morphizen-pattern/CMakeLists.txt: pass
--cpp_out=dllexport_decl=MORPHIZEN_PROTO_DLL_SPEC and define that macro on the
generated sources instead. Windows behavior is unchanged.

Test plan

Built and measured on an 8x MI350X host (gfx950, EPYC 9965, ROCm 7.2.0,
ONNX Runtime 1.25.1):

python build.py --build_dir /data/mm/onnxbuild/build_hipep \
  --install_dir /data/mm/onnxbuild/install_hipep --hip_arch gfx950 \
  --therock_dist /opt/rocm \
  --cmake_prefix_path /data/mm/onnxbuild/build/_deps/onnxruntime-src \
  --skip_submodule_sync --skip_wheel --skip_tests -j 384

Build is clean — no new warnings or errors.

Artifact-level equivalence (gfx950). The strongest evidence that this is inert
on CDNA:

Artifact	Pre-fix vs post-fix
`libhipgpu.so`	byte-identical (same SHA-256)
`libcustom_kernels_gfx950.so` `.hip_fatbin`	byte-identical — device ISA unchanged
`libcustom_kernels_gfx950.so` (whole file)	differs only in host launcher code

libhipgpu.so coming out bit-identical also validates the export split: every
morphizen TU and all six regenerated .pb.cc files recompiled against the new
header, and the link output did not change a byte. 2737 morphizen symbols remain
in .dynsym with GLOBAL DEFAULT binding.

End-to-end sweep. OGA model_benchmark, batch 1, generate 128, 1 warmup +
3 reps, via run_sweep.sh. See the Performance table below. Decode is flat within
1% and peak working set is identical everywhere.

Interleaved A/B. TTFT scatters in both directions across the sweep, so the
largest apparent gap (phi-4 L=128) was re-measured by swapping only
libcustom_kernels_gfx950.so and alternating arms A,B,A,B at 10 reps per launch:
pre averaged 48.5 ms (50.2 / 49.6 / 43.7 / 50.5), post averaged 46.0 ms
(45.0 / 44.4 / 48.5). The post-fix arm is marginally faster, so the sweep figure
was noise.

Output correctness. All six runs completed full 128-token generations with
coherent, on-topic text and no NaN/Inf tokens. Stricter text-equality is not
available on this harness: gpt-oss-20b has do_sample: true, top_k: 50, and even
greedy phi-4 is not reproducible run-to-run because the wave64 INT4 prefill path
benchmarks hipBLASLt candidates and caches the winner, so a timing-dependent
algorithm choice shifts rounding. Two runs of the same pre-fix binary already
diverge. Against that floor the post-fix build reproduces two of three greedy runs
exactly against each baseline — the same rate the two baselines achieve against
each other.

wave32 side. Verified by building for gfx1100 and diffing the generated device
ISA against the pre-wave64-port revision (unchanged apart from the content-hash
symbol). Not runtime-tested — see Notes.

Performance

Baseline is the pre-fix build measured on the same machine on the same day.

Metric	Model	L	Before	After
TTFT	phi-4	128	45.5 ms	53.9 ms
TTFT	phi-4	512	173.6 ms	141.7 ms
TTFT	phi-4	2048	550.1 ms	622.8 ms
TTFT	gpt-oss-20b	128	683.9 ms	711.2 ms
TTFT	gpt-oss-20b	512	2945 ms	3007 ms
TTFT	gpt-oss-20b	2048	11154 ms	10923 ms
Decode	phi-4	128	54.0 tok/s	54.1 tok/s
Decode	phi-4	512	106.5 tok/s	107.9 tok/s
Decode	phi-4	2048	64.6 tok/s	64.9 tok/s
Decode	gpt-oss-20b	128	238.8 tok/s	238.1 tok/s
Decode	gpt-oss-20b	512	210.5 tok/s	211.7 tok/s
Decode	gpt-oss-20b	2048	107.2 tok/s	110.4 tok/s

TTFT moves in both directions and is within run-to-run variance: within-run stddev
on phi-4 TTFT is 3–5 ms at L=128 and 42–54 ms at L=2048. The interleaved A/B above
settles the one case that looked like a regression.

Hardware, workload, build configuration, and measurement method: 8x AMD Instinct
MI350X (gfx950, CDNA4, wave64), pinned to device 0, others idle; AMD EPYC
9965 192-core host; ROCm 7.2.0; ONNX Runtime 1.25.1; models phi-4 (14B dense) and
gpt-oss-20b (MoE), both INT4 RTN block-32; OGA model_benchmark, batch 1,
generate 128, 1 warmup + 3 measured reps; build command as above.

github-actions · 2026-07-31T01:34:04Z

Thanks for opening a PR!

This project follows LLVM's incremental-development and AI-tool-use
guidance. See CONTRIBUTING.md
for the project workflow.

Before requesting review, please check that:

The change is focused. Substantial work links the relevant issue
or design discussion.
The PR documents relevant test results and updates affected
documentation.
If AI tools provided substantial assistance, the description
explains what was assisted and how it was validated, and commit
trailers identify the tool. The contributor has reviewed and
understands the result.

Reviewers are assigned through
CODEOWNERS where ownership
is configured.

The custom kernels were written for RDNA3/RDNA4 (wave32) around the WMMA matrix intrinsics and dot8-insts. CDNA runs wave64 and has neither, so the shared kernel sources did not compile and dispatch correctly there. Add lib/Runtime/Kernels/include/hip_arch_compat.h, a portability shim providing HIPDNN_WAVE_SIZE, HIPDNN_HAS_WMMA, a portable hipdnn_sudot4() (sudot4 on RDNA, the signed sdot4 on CDNA) and a runtime hipdnn_device_is_wave64() probe, and build on it: - gqa/qmoe/matmul_nbits kernels: make the warp-shuffle reductions span the real hardware wave and route the INT8 dp4a sites through the shim. - gqa.cpp: disable the WMMA-based fused and flash-decode paths on wave64, routing to the decomposed hipBLASLt pipeline. - matmul_nbits: guard the WMMA prefill fast paths off on wave64 and add a wave64 INT4 prefill path that dequantizes B to fp16 once (cached per weight pointer) and runs a hipBLASLt GEMM with a per-shape algo cache. - morphizen: leave MORPHIZEN_DLL_SPEC empty on non-MSVC. protobuf v34 emits `class <DLL_SPEC> [[gnu::warn_unused]] Name`, and GCC rejects a GNU __attribute__ placed before a C++11 [[...]] attribute. Validated on MI350X (gfx950, ROCm 7.2.0) with phi-4 and gpt-oss-20b INT4: both generate coherent output; phi-4 prefill sustains ~3.1-3.6k tok/s.

Disabling the fused path on wave64 makes fused_supported false there, which trips the quantized-KV guard added by #550 and fails with a message blaming fused_supported. That points a reader at the feature gates (causal, window, sink, bias) rather than the actual cause: the quantized-cache kernels exist only on the WMMA fused path, so a quantized cache has no implementation on CDNA at all. Reject the combination up front and name the architecture. Behaviour is unchanged -- both paths already returned -1 -- but the diagnostic now says what is wrong. RDNA is untouched: hipdnn_device_is_wave64() is false there.

The wave64 port raised the softmax block-size floor to a hard-coded 64 so a 32-thread block could not become a partial wave on CDNA and fold inactive lanes into the __shfl_xor reduction. That floor also applies on RDNA, where it forces a second wave for every row count <= 32 -- the common decode shape -- for no benefit. Take the floor from hipdnn_device_wave_size() instead, so it is 32 on wave32 (restoring the original geometry) and 64 on wave64. num_waves must come from the same value: it sizes the dynamic LDS fan-in buffer, which has to match the kernel's own blockDim.x / WAVE_SIZE. The host pass hardcodes HIPDNN_WAVE_SIZE to 32, so on gfx950 this also stops over-allocating that buffer by 2x. Also cache the wave-size probe per device ordinal rather than once per process. A host can drive several GPUs of different families (e.g. a gfx1151 APU beside a discrete card), where a single cached answer would be applied to the wrong device after hipSetDevice. The array is zero-initialized so no thread-safe-statics guard is emitted, and a failed query still yields the wave32 default. Verified on MI350X: the gfx950 device ISA is byte-identical before and after, and a benchmark sweep of phi-4 and gpt-oss-20b at L=128/512/2048 is unchanged within run-to-run variance.

Working around protobuf v34 by redefining MORPHIZEN_DLL_SPEC to nothing on non-Windows reached much further than the generated code it was meant for: that macro also decorates roughly thirty hand-written public API headers (graph.hpp, pass.hpp and friends). Nothing is built with -fvisibility=hidden today, so it is currently harmless, but it leaves a trap -- if hidden visibility ever reaches these targets, the public API silently stops being exported and only shows up as a link failure downstream. Introduce a separate MORPHIZEN_PROTO_DLL_SPEC for the protoc dllexport_decl and restore MORPHIZEN_DLL_SPEC to __attribute__((visibility("default"))). Only the new macro is empty on ELF, which is all the workaround needs: protobuf v34 emits `class <decl> [[gnu::warn_unused]] Name`, and GCC rejects a GNU __attribute__ placed before a C++11 [[...]] attribute. Windows behaviour is unchanged -- both macros still expand to __declspec(dllexport). Verified on MI350X: with every morphizen translation unit and all six generated .pb.cc files recompiled against the new header, libhipgpu.so links byte-identical to the previous build.

The wave64 enablement commits left gqa.cpp and matmul_nbits.cpp unformatted, so `lintrunner --apply-patches` rewrites them and the pre-commit hook fails with "files were modified by this hook". Both files are clean on main, so the drift is this branch's; these are the only two lint-dirty files in the tree. Pure reflow, no semantic change. The one exception is the hipBLASLt layout table in matmul_nbits.cpp: clang-format reflows that hand-aligned comment into unreadable soup, so it is fenced with clang-format off/on to keep the column alignment that makes it legible. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Made-with: Cursor

github-actions · 2026-07-31T02:09:47Z

L2 Accuracy Results (EP vs CPU)

Model	Combined L2	Total Elems
conv_test_hybrid	4.8668E-07	64
GroupQueryAttention_seq256	25.2366	2621440
MatMulNBits_o_seq128	259.906	368640
QMoE_seq128	34.957	368640

Threshold: 0.01 | Run: 3469 - Commit: 77c2ff0

github-actions · 2026-07-31T02:09:49Z

MorphiZen EP Performance Results

Model	QPS	Session (s)	1st Infer (ms)	CPU%	Mem (MB)
full_model_seq128	7.51	6.32	361	3	1242
GroupQueryAttention_seq128	4410.73	1.7547	11	6	311
matmul_down_seq128	514.80	2.46	89	3	354

EPContext Export Performance

Model	QPS	Session (s)	1st Infer (ms)	CPU%	Mem (MB)
full_model_seq128	7.52	46.65	361	3	15588

EPContext Import Performance

Model	QPS	Session (s)	1st Infer (ms)	CPU%	Mem (MB)
full_model_seq128	7.51	9.86	363	3	15758

OGA Benchmark Results

Model	Warmup	Reps	Prompt Len	Gen Tokens	TTFT (ms)	TPS	Peak Mem (GB)	GPU Mem (GB)
gpt-oss-20b-webgpu-int4-rtn-block-32	1	5	128	128	184.3	77.8	1.33	13.54
Llama-3.1-8B-awq-g128-int4-asym-fp16-onnx-dml	1	5	128	128	208.1	40.9	1.22	6.43

OGA Wheel Smoke (Python benchmark_e2e.py)

Model	TTFT (ms)	TPS
Llama-3.1-8B-awq-g128-int4-asym-fp16-onnx-dml	140	39.3

Run: 3469 - Commit: 77c2ff0

The wave64 enablement gated GQA's WMMA paths but missed MultiHeadAttention, which has its own fused flash-prefill kernel built on the same RDNA-only wave32 WMMA intrinsic. That kernel guards itself with __builtin_trap() on architectures without the builtin, on the assumption that the host never launches it -- but wrap_multi_head_attention had no wave-size check at all, so on CDNA the launch runs straight into the trap and aborts the queue: HSA_STATUS_ERROR_EXCEPTION: An HSAIL operation resulted in a hardware exception. code: 0x1016 Gate the fast path on !hipdnn_device_is_wave64(), mirroring fused_supported in gqa.cpp. The decomposed path below already handles every ineligible shape (decode Sq==1, causal/unidirectional, head-dim > 256), so wave64 simply takes it; RDNA is unaffected because the probe is false there. Found by running the E2E ctest suite on MI350X, which no CI job does -- linux-build.yml runs only the GPU-free LIT and unit targets. phi-4 and gpt-oss-20b use GQA rather than MHA, so no benchmark exercised this path. Verified on MI350X (gfx950, ROCm 7.2.0): E2E_Execute_test_multi_head_attention_model goes from aborting the queue to passing with output validation, and the rest of the suite is unchanged. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Made-with: Cursor

MI350X support landed on this branch but nothing documented it: docs/ mentions neither MI350, gfx950, CDNA nor wave64, and quick_start_linux.md assumes a gfx1151 target throughout. Add quick_start_mi350.md as a delta on the Linux guide rather than a third full copy. The build entry paths, $ROOT setup and the three benchmarking tools are identical on MI350X, so duplicating them would leave three files to keep in sync; the new guide links into those sections and covers only what differs -- why wave64 matters, the build and artifact specifics, what is disabled or unsupported on CDNA, the hipBLASLt dependency, the env knobs that are inert on wave64, measured numbers, and MI350X-specific troubleshooting. Every claim was reproduced on an 8x MI350X host: build.py auto-detects gfx950 from gfx_target_version 90500; a cold build produces libhipgpu.so alongside libcustom_kernels_gfx950.so; a missing or wrong-arch kernel library is a hard failure at JIT init, which is why the published CI package (gfx1151 only) cannot run here; and the performance table reproduces on an independently built tree within the run-to-run variance the guide documents. Linux only -- MI350X is not supported on the Windows flow. Also link the new guide from README.md and quick_start_linux.md. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Made-with: Cursor

penglu12-quark requested review from ChaoLi-AMD, jiangyon-amd and qianglin-amd July 31, 2026 01:33

penglu12-quark self-assigned this Jul 31, 2026

penglu12-quark requested review from amd-mingw, edelaye, fhanuman and wcy123 as code owners July 31, 2026 01:33

penglu12-quark added the enhancement New feature or request label Jul 31, 2026

penglu12-quark and others added 5 commits July 30, 2026 21:39

penglu12-quark force-pushed the users/penglu12/enablement_for_mi350 branch from 0e1d6bb to 40fc3b6 Compare July 31, 2026 01:44

penglu12-quark and others added 2 commits July 31, 2026 02:58

ChaoLi-AMD changed the title ~~feat/penglu12/enablement for mi350~~ HIP EP Enablement for MI350 Jul 31, 2026

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

HIP EP Enablement for MI350 - #602

HIP EP Enablement for MI350#602
penglu12-quark wants to merge 7 commits into
mainfrom
users/penglu12/enablement_for_mi350

penglu12-quark commented Jul 31, 2026

Uh oh!

github-actions Bot commented Jul 31, 2026

Uh oh!

github-actions Bot commented Jul 31, 2026 •

edited

Loading

Uh oh!

github-actions Bot commented Jul 31, 2026 •

edited

Loading

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

Uh oh!

Conversation

penglu12-quark commented Jul 31, 2026

Summary

Why

What

Test plan

Performance

Uh oh!

github-actions Bot commented Jul 31, 2026

Uh oh!

github-actions Bot commented Jul 31, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

L2 Accuracy Results (EP vs CPU)

Uh oh!

github-actions Bot commented Jul 31, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

MorphiZen EP Performance Results

EPContext Export Performance

EPContext Import Performance

OGA Benchmark Results

OGA Wheel Smoke (Python benchmark_e2e.py)

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

github-actions Bot commented Jul 31, 2026 •

edited

Loading

github-actions Bot commented Jul 31, 2026 •

edited

Loading