HIP backend: kernels taking `__bf16` never launch on ROCm 7.1 / clang 21 (host stub and device symbol mangled differently)

# HIP backend: kernels taking `__bf16` never launch on ROCm 7.1 / clang 21 (host stub and device symbol mangled differently)

**Environment**

| | |
|---|---|
| comfy-kitchen | 0.2.33 (built from source: `comfy_kitchen/backends/hip`) |
| GPU / arch | AMD Radeon AI PRO R9700 — `gfx1201` (RDNA4, wave32) |
| ROCm | 7.1 (`/opt/rocm`, `libhsa-runtime64.so.1.18.0`, HIP 7.1.52801) |
| Compiler | clang **21.1.8** (the ROCm/Ubuntu toolchain, `/opt/rocm/llvm/bin/clang++`) |
| OS / arch | Ubuntu, aarch64 |
| Build cmd | `cmake -S comfy_kitchen/backends/hip -B build -G Ninja -DCMAKE_BUILD_TYPE=Release -DCOMFY_HIP_ARCHS=gfx1201 -Dnanobind_DIR=/nanobind/cmake -DPython_EXECUTABLE=/bin/python` |

**Symptom**

The extension builds and loads fine (`hip` backend reports `available: True`, 39 capabilities), and kernels that don't touch `__bf16` work (`quantize_int8_rowwise` ✅). But **every kernel whose signature mentions `__bf16` fails at launch** with:

```
:0:./hipamd/src/hip_global.cpp:109 : us: Cannot find Symbol with name:
_ZN5comfy11hip_backend16gemm_wmma_kernelINS0_7MmaInt8ENS0_10EpiRowwiseEu6__bf16Li64ELi64ELi64ELi2ELi2ELi2ELi2EEEvPKhS5_PT1_iiiiT0_
```

Affected: all `gemm_wmma_kernel<..., __bf16, ...>` instantiations (int8 / int4 / fp8 GEMMs), `adaln`, `rms_adaln`, `dequantize_per_tensor_fp8`, `convrot_w4a4_linear`, `gemv_awq`, the rope kernels and the sage/int8-attention family — i.e. anything instantiated with a bf16 output type or taking a `__bf16*` parameter. The official test suite cannot run: `pytest tests/test_hip_wmma.py` → 399 skipped ("HIP extension not built") once the module path is fixed, or the process aborts (`Fatal Python error`) when the int8 cases actually launch.

**Root cause: `__bf16` is mangled differently in the host pass and in the device pass**

Minimal reproduction (no comfy-kitchen needed) — `repro/bf16-mangling-repro.sh` in this report:

```cpp
#include
#include
__global__ void k_builtin(__bf16* p) { if (threadIdx.x == 9999) p[0] = __bf16(1.0f); }
__global__ void k_hipbf16(hip_bfloat16* p) { if (threadIdx.x == 9999) p[0] = hip_bfloat16(1.0f); }
void h(__bf16* a, hip_bfloat16* b) { k_builtin<<<1,1>>>(a); k_hipbf16<<<1,1>>>(b); }
```

```
$ hipcc --offload-arch=gfx1201 -c repro.hip -o repro.o
$ strings repro.o | grep -oE "_Z[0-9]+k_(builtin|hipbf16)[^ ]*" | sort -u # host-side launch stub names
_Z9k_builtinPu6__bf16 ← host pass mangles __bf16* as u6__bf16
_Z9k_hipbf16P12hip_bfloat16 ← hip_bfloat16* is consistent
$ # device-side symbol table:
_Z9k_builtinPDF16b ← device pass mangles the same __bf16* as DF16b ← MISMATCH
_Z9k_hipbf16P12hip_bfloat16 ← consistent
```

Same for template arguments:

```
host: _Z6k_tmplIu6__bf16EvPT_
device: _Z6k_tmplI12hip_bfloat16EvPT_ (for the same instantiation)
```

So the host-side launch stub embeds `...u6__bf16...` while the device code object exports `...DF16b...`; the HIP runtime looks the symbol up by the host-side name and never finds it. Note the `__AMDGCN_WAVEFRONT_SIZE__` deprecation warnings in clang 21 suggest the toolchain has been changing builtin handling in this cycle.

**Evidence from a real build (comfy-kitchen HIP extension)**

Diffing the host-referenced names against the device code objects extracted from the same `.so`
(`llvm-objcopy --dump-section .hip_fatbin=…`):

```
host-referenced MmaInt8 instantiations: 20 (includes EpiRowwise + u6__bf16 + 5 tile variants)
device-symbol MmaInt8 instantiations: 15 (same, but only 6__half / DF16b / f outputs)
difference (host has, device lacks): exactly the 5 ...Eu6__bf16... instantiations
reverse difference (device has, host lacks): empty
```

i.e. the only thing missing from the device code is the bf16-output spelling.

**What does NOT fix it**

- `-mllvm -amdgpu-internalize-symbols=false` (verified in the compile command line, still fails — so it's not internalization)
- `-fclang-abi-compat=17` / `=18`
- Rebuilding does emit the kernels (all 15 `ops/*.hip` TUs compile; `adaln.hip.o`, `gemm_int8.hip.o`, `per_tensor_fp8.hip.o` are all in the build dir) — the symbols are simply stamped with the device-pass spelling.

**Workaround that works (fully verified)**

Spell the type as `hip_bfloat16` in every *signature position* (kernel parameter types, template arguments at launch sites, and the dtype-dispatch macro invocations), while leaving `__bf16` everywhere it is a hardware vector type (`ext_vector_type`, `mma.h`'s `Elem`, `v8bf/v16bf`).

Three passes are needed; the third is the easy one to miss:

1. kernel definitions / launches — lines containing `__global__`, `<<<`, `launch_*<`, `store_out<`, `rope_store<`, `operator()<`;
2. the remaining parameter types in the same file (otherwise helper ↔ caller types stop matching);
3. **template arguments inside the dtype-dispatch macros** — `COMFY_LAUNCH_QUANT(__bf16)`, `COMFY_LAUNCH_DEQUANT(__bf16)`, `CK_LAUNCH_DEQUANT_CONVROT(__bf16)`, `CK_ROPE_LAUNCH(__bf16)`, `CK_RMS_ROPE_LAUNCH(__bf16)`, `CK_AWQ_LAUNCH(__bf16)`, `CK_SVD_LAUNCH(__bf16,…)`, `COMFY_REQUANT_LAUNCH(__bf16)`, `SAGE_QK_DISPATCH_DTYPE(__bf16)`. Missing these leaves the corresponding op broken even after 1 and 2.

After the change (and adding `#include ` to the touched files):

```
$ cd /tmp && python -m pytest /tests/test_hip_wmma.py /tests/test_hip_dispatch.py -q
451 passed, 2 skipped, 1 warning in 20.62s (was: 399 skipped + process abort)
```

Numerical cross-check against the `eager` backend for the previously-crashing ops: `int8_linear` rel. mean err 7.7e-3; `adaln`, `rms_adaln`, `dequantize_per_tensor_fp8` bit-identical. Performance on gfx1201: int8 GEMM 4096³ = **155 TFLOPS** (vs 34 TFLOPS for the eager fallback), `sol_attn`/`int8_attention` run as WMMA kernels (int8 attention ~100 TFLOPS at S=4096·H=16·D=128 vs 48 TFLOPS for SDPA).

**Side note (separate, minor): `setup.py --hip` cannot configure on this toolchain**

`setup.py`'s CMake invocation passes `-DCMAKE_C_COMPILER=/opt/rocm/bin/hipcc -DCMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc`; with hipcc as the C/CXX compiler CMake configuration fails. Using default system compilers (gcc 15 for C/CXX, clang 21 discovered for HIP) configures fine. Also, if `COMFY_HIP_ARCHS` is not passed, the build targets all 17 archs.

**Suggestion**

Mangle-stable handling would be to avoid the builtin `__bf16` in mangled positions (use `hip_bfloat16`, or a fixed-width carrier such as `unsigned short` with a reinterpret_cast at kernel entry), and/or use an alias that both passes mangle identically. A `static_assert`-style check at build time (compile a small TU that launches a `__bf16` kernel and compare host stub vs device symbol) would catch this class of breakage early on new toolchains.

Happy to provide the patched tree / full write-up if useful.

Start with the __bf16 signature and dispatch-macro occurrences in the ops/*.hip files, then inspect repro/bf16-mangling-repro.sh to confirm the host/device symbol mismatch. Run tests/test_hip_wmma.py and tests/test_hip_dispatch.py on ROCm 7.1; done means the affected kernels launch, the tests pass, and the reproduced symbols use matching names.

Written by the indexing model from the issue text.