Comfy-Org/comfy-kitchen

Fork 0

mirror of https://github.com/Comfy-Org/comfy-kitchen.git synced 2026-11-04 23:37:42 +00:00

[GH-ISSUE #170] HIP backend: kernels taking `__bf16` never launch on ROCm 7.1 / clang 21 (host stub and device symbol mangled differently) #54

New issue

Open

opened 2026-10-26 22:19:44 +00:00 by bsakaguchi · 1 comment

bsakaguchi commented

2026-10-26 22:19:44 +00:00

Owner

Originally created by @dyc-dev on GitHub (Sep 10, 2026).
Original GitHub issue: https://github.com/Comfy-Org/comfy-kitchen/issues/170

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=<venv>/nanobind/cmake -DPython_EXECUTABLE=<venv>/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 : <ts> 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:

#include <hip/hip_runtime.h>
#include <hip/hip_bfloat16.h>
__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:

kernel definitions / launches — lines containing __global__, <<<, launch_*<, store_out<, rope_store<, operator()<;
the remaining parameter types in the same file (otherwise helper ↔ caller types stop matching);
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 <hip/hip_bfloat16.h> to the touched files):

$ cd /tmp && python -m pytest <repo>/tests/test_hip_wmma.py <repo>/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.

Originally created by @dyc-dev on GitHub (Sep 10, 2026). Original GitHub issue: https://github.com/Comfy-Org/comfy-kitchen/issues/170 # 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=<venv>/nanobind/cmake -DPython_EXECUTABLE=<venv>/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 : <ts> 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 <hip/hip_runtime.h> #include <hip/hip_bfloat16.h> __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 <hip/hip_bfloat16.h>` to the touched files): ``` $ cd /tmp && python -m pytest <repo>/tests/test_hip_wmma.py <repo>/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.

bsakaguchi commented

2026-10-26 22:19:45 +00:00

Author

Owner

@0xDELUXA commented on GitHub (Sep 12, 2026):

Confirmed the diagnosis. With stock clang 21.1.6, AArch64 mangles __bf16 as u6__bf16 on the host and DF16b on the AMDGPU device, so the stub registers a symbol the code object does not export. x86_64 uses DF16b on both sides, which is why this is AArch64-only.

This is an upstream, deliberate ABI rule: AArch64/ARM retain u6__bf16. llvm/llvm-project#115521 and llvm/llvm-project#115956 are still open, and upgrading clang/ROCm does not fix it. AMD also does not support AArch64 hosts for ROCm, so this only affects unsupported, source-built configurations.

Given that, a broad change to avoid __bf16 in all mangled positions is not worth carrying for now. The local workaround should stay: avoid __bf16 in kernel parameters, template arguments, and dtype-dispatch macros; use it only for vector element types and casts inside kernel bodies. Prefer __hip_bfloat16 from <hip/hip_bf16.h> over the older hip_bfloat16.

The architecture fallback is deliberate. setup.py first honors an explicit override, otherwise uses visible AMD GPUs, and only falls back to the full DEFAULT_HIP_ARCHS list when no GPU can be enumerated. CMakeLists.txt follows the same logic. Getting all 17 means no GPU was detected at build time, so set COMFY_HIP_ARCHS=gfx1201 to build only that target.

The setup.py --hip failure with hipcc as CC/CXX is unrelated and should be tracked separately with the CMake configure log.

@0xDELUXA commented on GitHub (Sep 12, 2026): Confirmed the diagnosis. With stock clang 21.1.6, AArch64 mangles `__bf16` as `u6__bf16` on the host and `DF16b` on the AMDGPU device, so the stub registers a symbol the code object does not export. x86_64 uses `DF16b` on both sides, which is why this is AArch64-only. This is an upstream, deliberate ABI rule: AArch64/ARM retain `u6__bf16`. [llvm/llvm-project#115521](https://github.com/llvm/llvm-project/issues/115521) and [llvm/llvm-project#115956](https://github.com/llvm/llvm-project/pull/115956) are still open, and upgrading clang/ROCm does not fix it. AMD also does not support AArch64 hosts for ROCm, so this only affects unsupported, source-built configurations. Given that, a broad change to avoid `__bf16` in all mangled positions is not worth carrying for now. The local workaround should stay: avoid `__bf16` in kernel parameters, template arguments, and dtype-dispatch macros; use it only for vector element types and casts inside kernel bodies. Prefer `__hip_bfloat16` from `<hip/hip_bf16.h>` over the older `hip_bfloat16`. The architecture fallback is deliberate. `setup.py` first honors an explicit override, otherwise uses visible AMD GPUs, and only falls back to the full `DEFAULT_HIP_ARCHS` list when no GPU can be enumerated. `CMakeLists.txt` follows the same logic. Getting all 17 means no GPU was detected at build time, so set `COMFY_HIP_ARCHS=gfx1201` to build only that target. The `setup.py --hip` failure with `hipcc` as `CC/CXX` is unrelated and should be tracked separately with the CMake configure log.

bsakaguchi referenced this issue

2026-10-26 22:20:28 +00:00

[PR #54] Change rounding method in quantization process (triton backend fix for amd gpus) #109

No labels

pull-request

No milestone

No project

No assignees

1 participant

Notifications

Due date

The due date is invalid or out of range. Please use the format "yyyy-mm-dd".

No due date set.

Dependencies

No dependencies set.

Reference

Comfy-Org/comfy-kitchen#54

No description provided.

Rows
Columns

[GH-ISSUE #170] HIP backend: kernels taking __bf16 never launch on ROCm 7.1 / clang 21 (host stub and device symbol mangled differently) #54

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

[GH-ISSUE #170] HIP backend: kernels taking `__bf16` never launch on ROCm 7.1 / clang 21 (host stub and device symbol mangled differently) #54

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