ROCm 7 build fails: #3267 compat probes test for macros that HIP 7 delivers as enum members and functions

[KhazAkar]

Environment

• OS: Fedora 44
• ROCm 7.1.52802 (HIP_VERSION_MAJOR = 7), clang from the ROCm toolchain
• Targets: gfx1102 (Radeon RX 7700S 8GB VRAM) and gfx1103 (Radeon 780M, UMA)
• Host: Ryzen 9 7940HS, Framework 16, 64 GB DDR5 RAM, ~54GB usable for the host
• Tree: main after rocm: fix ROCm 5.7 / clang-17 build on Strix (gfx1151) #3267, cmake HIP preset, tests excluded

Problem

The ROCm 5.7 compatibility layer from #3267 uses #ifndef probes. Two of
them misfire on ROCm 7, and the build fails for both RDNA3 targets.

1. src/vt/rocm/rocm_matmul_hipblaslt.hip probes #ifndef HIPBLAS_COMPUTE_32F.
   On ROCm 7 that name is an enumerator of hipblasComputeType_t in
   hipblas-common.h, not a macro. The probe reads true, so the pre-ROCm6
   block activates and typedefs hipblasDatatype_t, a type ROCm 7 removed.
   Result: 20 compile errors in that file.

2. include/vt/rocm/hip_shfl_compat.h probes #ifndef __shfl_down_sync and
   #ifndef __shfl_sync. On ROCm 6 and later those names exist as functions
   in amd_warp_sync_functions.h, invisible to #ifndef. The macro shims
   activate anyway and collide with the real declarations.

Root cause

#ifndef answers "does a macro with this name exist". HIP delivers API
changes as enumerators and functions, so the probe tests the wrong property.
Every ROCm release that moves a name from macro to enum or function form
silently re-arms these branches.

Proposed fix

Gate version-dependent shims on HIP_VERSION_MAJOR, and keep #ifndef
probes only for names that stay macros across releases:

// hip_shfl_compat.h: the _sync wrappers exist as functions on ROCm 6+
#if !defined(HIP_VERSION_MAJOR) || HIP_VERSION_MAJOR < 6
...existing shims...
#endif

// rocm_matmul_hipblaslt.hip: enum form landed in ROCm 6,
// hipblasDatatype_t was removed in ROCm 7
#if !defined(HIPBLAS_COMPUTE_32F) && \
    (!defined(HIP_VERSION_MAJOR) || HIP_VERSION_MAJOR < 7)
...existing typedef and aliases...
#endif

Other propositions

Maybe... we should also think about shared header for ROCm versioning handling?

Other notes

I already have local version with fix. Tested on Granite 4.2 3b safetensors, fp8 and bf16. Results on iGPU and dGPU:

Device	bf16	fp8
7700S	22.8 t/s, 6.9G	22.3 t/s, 3.9G
780M	10.2 t/s	8.8 t/s