Skip to content

TODO.md

Latest commit

 

History

History
26 lines (16 loc) · 3.5 KB

TODO.md

File metadata and controls

26 lines (16 loc) · 3.5 KB
Raw

ROCm Backend — Roadmap

中文版 / Chinese: TODO.zh-CN.md

The AMD ROCm/HIP backend implements the same acl* runtime + aclnn* operator + Hccl* contract as the CUDA and Metal backends, so the backend-agnostic clients in ../tests/ and ../tools/ link against it without source changes (make -C ../tests BACKEND=../rocm).

Operator coverage is functionally complete on this machine — all 42 backend-agnostic test binaries pass (make -C ../tests BACKEND=../rocm run → RUN_EXIT=0), and Qwen3-0.6B runs end-to-end matching HuggingFace (prefill logits 5.582e-6, greedy 8/8). What this file tracks is what is left or bound to hardware this machine does not have: the low-precision matrix-core path that RDNA3.5 lacks, multi-rank collectives, bit-exact NZ/fractal fidelity, and build/perf headroom.

Scope mirrors the project as a whole: functional / tolerance-level validation and cross-architecture backend exploration — not bit-exact Ascend fidelity, not a cross-vendor performance contest.

Target: gfx1151 (Strix Halo APU, RDNA3.5, wavefront size 32), ROCm 7.2.4, unified LPDDR5X (~256 GB/s theoretical, shared with the CPU).

A. Bound to hardware this machine does not have (RDNA3.5 / single APU)

  • Native low-precision matrix cores. fp8 / MXFP8 / MXFP4 GEMM currently run the fp16-dequant functional fallback, and CANN_FAST_TF32 is a no-op, because RDNA3.5 has no fp8 / microscale / TF32 matrix cores. mx_use_hw() is gated off under __HIP_PLATFORM_AMD__; a native microscale GEMM would light up on CDNA (gfx942/gfx950) via MXFP8_FORCE_HW=1, but validating it needs a CDNA machine. On gfx1151 the functional path is both correct and the only path with a payoff (fp8 is 2–3× slower than fp16 here — memory-footprint tier only; see BENCHMARK.md).
  • Multi-GPU / RCCL scaling. RCCL collectives are single-rank-validated only (one APU). test_dist_parity (the 2-node RoCE test) is N/A here, and the mc2 fused-collective + multi-rank paths are untested on ROCm hardware. A genuine multi-GPU / multi-node ROCm box (RoCE or Infinity Fabric) is needed to exercise the non-degenerate collective paths.
  • Bit-exact NZ / fractal-layout fidelity. The Ascend-specific fused / quant / distributed / NZ-format ops are verified by structural invariants, degenerate-equivalence, and the shared tolerance suite. Confirming the NZ fractal layout matches Ascend's native layout bit-for-bit needs real Ascend hardware — the same permanent out-of-scope item the CUDA and Metal backends carry.

B. Performance / build headroom (not correctness)

  • MIOpen algorithm selection on this build. miopenFindConvolutionForwardAlgorithm returns 0 candidates on this ROCm/gfx1151 build (a safe ImplicitGEMM default is used); revisit on a newer MIOpen. First-conv-per-shape latency is ~5–9 s (Find + kernel compile), amortized by the algo cache — a persistent on-disk cache could remove the cold-start cost across process runs.
  • FlashAttention-3-grade attention kernel. The rocWMMA flash path (16×16×16, wave32) mirrors the CUDA WMMA kernel and carries the same occupancy headroom; the hipBLASLt batched perf-path is the default when S/P materialization fits. An FA3-class redesign (smaller shared-mem tiles, warp specialization) is the same open item the CUDA backend tracks.
  • Incremental build time. The build uses -fgpu-rdc whole-program device link. Exploring per-TU compile + a separate device-link step would speed up incremental rebuilds.