RDNA2 issues on Linux (ROCm)

[tacuna]

I had errors while training/building on RDNA 2, 6700 XT (gfx1030)the using ROCM.
(this issue has been identified by AI, i confirmed the fixes work at least for my GPU. I'm not knowledgeable enough myself to judge to correctness of the fixes, hence did not create a MR)

Environment

• GPU: AMD Radeon RX 6700 XT (gfx1031, HSA_OVERRIDE_GFX_VERSION=10.3.0)
• ROCm: 7.13.0 (Arch Linux rocm-gfx103x-bin)
• CubeCL: 0.11.0-pre.1 (rev a8ddb49dc928d0ea6cfc86101013082e24a5ff18)
• Burn: 0.22.0-pre.1 (via burn-rocm)
• Rust: nightly 1.98

Summary

On ROCm/HIP (AMD GPU), any kernel that uses a vectorized float type (e.g. float × 4) fails to
compile at runtime with errors like:

error: unknown type name 'float_4'; did you mean 'float4'?
   struct __attribute__((aligned(16))) array<float_4, 1> {
                                             ^~~~~~~
                                             float4

error: template specialization requires 'template<>'
   struct __attribute__((aligned(16))) array<float_4, 1> {

error: expected unqualified-id
   struct __attribute__((aligned(16))) float_4* {

Root cause

In crates/cubecl-cpp/src/shared/kernel.rs, type_vectorized_definitions iterates a
HashSet<Item<D>> and generates a struct definition for every item whose
vectorization() > 1. However, Item::Pointer and Item::Array inherit their vectorization
count from their inner Item::Vector, so they also match the > 1 condition:

Item variant	Display output	vectorization()
Item::Vector(float, 4)	float_4	4
Item::Pointer(float_4, Global)	float_4*	4 (inherited)
Item::Array(float_4, 1)	array<float_4, 1>	4 (inherited)

This causes two separate problems:

1. Invalid struct names — float_4* and array<float_4, 1> are not valid C++ struct
   identifiers (a pointer type and a template specialization cannot be the subject of a struct
   definition body). hiprtc rejects them with parse errors.

2. Use-before-definition — HashSet iteration order is non-deterministic. When
   array<float_4, 1> is emitted before the float_4 struct that it depends on, hiprtc
   sees an unknown type name on the very first use.

The generated (broken) preamble looks like:

// HashSet iteration gave this order (undefined, varies per run):
struct __align__(16) array<float_4, 1> {  // ← float_4 not yet defined!
    float i_0; float i_1; float i_2; float i_3;
};
struct __align__(16) float_4* {           // ← '*' in struct name is invalid C++
    float i_0; float i_1; float i_2; float i_3;
};
struct __align__(16) float_4 {            // ← defined too late
    float i_0; float i_1; float i_2; float i_3;
};

Fix

In type_vectorized_definitions:

1. Filter to only Item::Vector — the only variant that actually needs its own struct definition.
   Item::Pointer and Item::Array reference the vector struct by name at the point of use and do
   not need (and must not have) their own struct bodies.

2. Sort by (elem.size(), vectorization) before emitting, so that smaller/simpler types are
   always defined before larger ones. This gives a stable, forward-declaration-safe order.

pub fn type_vectorized_definitions<D: Dialect>(
    f: &mut std::fmt::Formatter<'_>,
    items: &HashSet<Item<D>>,
) -> std::fmt::Result {
    let mut vec_items: Vec<_> = items
        .iter()
        .filter(|item| matches!(item, Item::Vector(..)))
        .collect();
    vec_items.sort_by_key(|item| (item.elem().size(), item.vectorization()));

    for item in vec_items {
        let elem = item.elem();
        let size = item.vectorization();
        let alignment = elem.size() * size;
        if size > 1 {
            write!(f, "\nstruct __align__({alignment}) {item} {{")?;
            for i in 0..size {
                write!(f, "\n    {elem} i_{i};")?;
            }
            f.write_str("\n};")?;
        }
    }
    Ok(())
}

Reproduction

Any kernel that uses a float × 4 (or similar vectorized) item in a context that also produces
Pointer or Array wrapping that vector type — for example, a matmul with tensor shapes
m=7, n=611, k=611 in F32 on a HIP backend.

The crash manifests as an autotune panic:

thread 'main' panicked at cubecl-runtime/src/tune/tuner.rs:210:
Can't execute the autotune plan for key: MatmulAutotuneKey { definition:
  MatmulProblemDefinition { m: 7, n: 611, k: 611, ... elem_lhs: Scalar(Float(F32)), ... } }
 - results: [AutotuneResult { outcome: Err(matmul_naive: An unknown error happened.
     Caused by: A compilation error happened during launch
     Caused by: [Compilation Error]
       /tmp/.../input/CompileSource...:19:43: error: unknown type name 'float_4'

Note that all autotune candidates fail (including matmul_naive, which is supposed to always
work), because the compilation error occurs in the shared kernel preamble before any kernel-specific
code is reached.

---

Additional bug: is_wmma_capable returns true for GFX10 (RDNA2)

File

crates/cubecl-cpp/src/hip/arch.rs

Root cause

AMDArchitecture::is_wmma_capable includes GFX10 in its match arm:

// upstream (broken):
fn is_wmma_capable(&self) -> bool {
    matches!(self, AMDArchitecture::GFX10 | AMDArchitecture::GFX11 | AMDArchitecture::GFX12)
}

RDNA2 (GFX10 / gfx1030–gfx1032) does not support AMD WMMA intrinsics. WMMA is only
available on RDNA3 (GFX11) and RDNA4 (GFX12). When is_wmma_capable returns true for
GFX10, WmmaIntrinsicCompiler::supported_wmma_combinations registers cooperative-matrix
matmul kernels that all fail at runtime, causing the autotune benchmark for those kernels
to panic and mark every candidate as invalid.

Fix

fn is_wmma_capable(&self) -> bool {
    matches!(self, AMDArchitecture::GFX11 | AMDArchitecture::GFX12)
}

---

Additional bug: RocWmmaCompiler::supported_wmma_combinations includes GFX10

File

crates/cubecl-cpp/src/hip/mma/rocwmma_compiler.rs

Root cause

The supported_wmma_combinations method in RocWmmaCompiler (the compiler used when the
rocwmma feature is enabled) has a combined match arm for GFX10 | GFX11:

// upstream (broken):
AMDArchitecture::GFX10 | AMDArchitecture::GFX11 => {
    // ... returns non-empty list of WMMA combinations
}

This registers WMMA combinations for GFX10 even though RDNA2 has no WMMA support,
causing the same crash described above when the rocwmma feature is enabled.

Fix

Split the arm and return an empty list for GFX10:

AMDArchitecture::GFX10 => vec![],  // RDNA2 has no WMMA support
AMDArchitecture::GFX11 => {
    // ... existing GFX11 logic unchanged
}

---

Related: hipGetDevicePropertiesR0600 — versioned symbol in cubecl-hip-sys

What it is

It lives in the upstream cubecl-hip-sys
crate (crates.io, version 7.1.5280200), which provides the raw bindgen-generated FFI bindings
for libamdhip64.so.

In ROCm 6.0, AMD versioned the hipDeviceProp_t struct and introduced a renamed symbol
hipGetDevicePropertiesR0600 (the _R0600 suffix encodes the struct ABI version). Older ROCm
(< 6.0) exported only the plain hipGetDeviceProperties. cubecl-hip-sys 7.x targets ROCm 6+
and therefore links against hipGetDevicePropertiesR0600.

Why it breaks after a system update

If the system update downgrades or partially replaces the HIP runtime (e.g. libamdhip64.so),
the versioned symbol may disappear, causing a link error:

error: linking with `cc` failed: exit status: 1
  = note: /usr/bin/ld: cannot find -lamdhip64
  = note: undefined reference to `hipGetDevicePropertiesR0600`

Check:

nm -D /opt/rocm/lib/libamdhip64.so | grep hipGetDeviceProperties
# Should show: hipGetDevicePropertiesR0600 (ROCm 6+)
#              hipGetDeviceProperties      (ROCm 5 and earlier, or compatibility alias)
cat /opt/rocm/.info/version     # or: hipconfig --version

If the symbol is missing, the ROCm packages are in a mixed state. Fix:

# Arch Linux (using rocm-gfx103x-bin — do NOT install hip-runtime-amd, it conflicts)
sudo pacman -Syu
sudo pacman -S --reinstall rocm-gfx103x-bin

# Or set HIP_PATH to a known-good install and let cargo rebuild:
export HIP_PATH=/opt/rocm
cargo clean
cargo build --features rocm

Could the versioned symbol be detected dynamically?

Yes, in principle — and cubecl-hip-sys already does this correctly. Its build.rs
runs hipconfig --version at compile time to determine the installed HIP patch version and
selects the matching bindings_XXXXX.rs module via a Cargo feature (hip_52802, hip_51831,
etc.). The _R0600 binding is only activated for the ROCm 6.0+ feature flags.

So the mechanism is dynamic (build-time detection), not hardcoded. The link failure is a
runtime/installation problem (wrong .so on disk after a partial upgrade), not a flaw in
the version-selection logic. The fix is always to restore a consistent ROCm installation.

cubecl-hip-sys version in use

Field	Value
Crate	cubecl-hip-sys
Version	7.1.5280200 (HIP patch 52802, ROCm 7.1.1)
Source	crates.io
Requires	ROCm 6.0+ (libamdhip64.so must export hipGetDevicePropertiesR0600)

Arch Linux workaround (hipconfig wrapper)

On Arch with rocm-gfx103x-bin the hipconfig binary has three problems that break the
cubecl-hip-sys build script (panics on stderr, ANSI codes in version string, unknown patch
version 99004).

1. hipconfig writes warnings to stderr (unofficial gfx1031 hardware). The build script
   panics on any stderr output → E0425 cubecl_hip_sys not found / E0432.
2. hipconfig --version embeds ANSI escape codes inside the version number string,
   corrupting the patch-number regex parse → panic.
3. The reported patch version (99004 from 7.13.99004-hash) is beyond the range known
   to cubecl-hip-sys (max 52802 for ROCm 7.1) → no bindings selected → E0425.

Fix: replace hipconfig with a wrapper that returns hardcoded clean values:

 mkdir -p ~/.local/bin
 printf '#!/bin/sh\n[ "$1" = "--version" ] && echo "7.1.52802-fake" && exit\n[ "$1" = "-R" ] && echo "/opt/rocm" && exit\nexec /opt/rocm/bin/hipconfig "$@" 2>/dev/null\n' \
     > ~/.local/bin/hipconfig
 chmod +x ~/.local/bin/hipconfig
 # Add to ~/.bashrc so it persists:
 echo 'export PATH=~/.local/bin:$PATH' >> ~/.bashrc
 export PATH=~/.local/bin:$PATH

The wrapper returns 7.1.52802-fake for --version (selects hip_52802 ROCm 7.1
bindings, ABI-compatible with 7.13) and /opt/rocm for -R (library search path).
Both are clean stdout with no stderr and no ANSI codes.