hip/gfx906: cudaFuncSetAttribute must be a function, not a macro (ROCm 7.2.1 build fix)

[xxDoman]

Problem

Building the gfx906 backend (-DSTRATA_HIP_GFX906=ON) from source fails on ROCm 7.2.1: cudaFuncSetAttribute is a function-like macro in include/strata/platform/hip_compat/strata_hip.h, and the call sites pass a template instantiation whose argument list contains commas — the preprocessor splits on them and the macro gets too many arguments.

Environment

OS	Ubuntu 26.04.1 LTS, kernel 7.0.0-38, x86_64
CPU	Intel Core i5-12400F
GPU	AMD Instinct MI50 32 GB — gfx906 (Vega 20, wave64)
ROCm	7.2.1
HIP compiler	AMD clang 22.0.0git (roc-7.2.1, commit 26084)
CMake / Make	4.2.3 / 4.4.1 · gcc 15.2.0 · glibc 2.43

cmake -DSTRATA_HIP_GFX906=ON -DCMAKE_HIP_ARCHITECTURES=gfx906 \
      -DCMAKE_HIP_COMPILER=/opt/rocm-7.2.1/lib/llvm/bin/clang++ \
      -DCMAKE_BUILD_TYPE=Release ..
make -j11 strata      # EXIT 2

Error

include/strata/platform/hip_compat/strata_hip.h:154:

#define cudaFuncSetAttribute(fn, attr, val) hipFuncSetAttribute(reinterpret_cast<const void*>(fn), attr, val)

src/kernels/cuda/fused_gr.cu:1089-1093:

cudaFuncSetAttribute(gr_down_v3_kernel<1, kFusedGrMaxT, false>, cudaFuncAttributeMaxDynamicSharedMemorySize, need1) == cudaSuccess &&
cudaFuncSetAttribute(gr_down_v3_kernel<1, 6, true>, cudaFuncAttributeMaxDynamicSharedMemorySize, need6) == cudaSuccess) {
    cudaFuncSetAttribute(gr_down_v3_kernel<1, 4, true>, cudaFuncAttributePreferredSharedMemoryCarveout, 100);
    ...

fused_gr.cu:1089:81: error: too many arguments provided to function-like macro invocation
strata_hip.h:154:9: note: macro 'cudaFuncSetAttribute' defined here
fused_gr.cu:1089:17: error: use of undeclared identifier 'cudaFuncSetAttribute'

5 errors (lines 1089-1093). The preprocessor reads gr_down_v3_kernel<1, kFusedGrMaxT, false> as several macro arguments (<1, kFusedGrMaxT, …) and chokes on the rest. This is standard preprocessor behaviour, not a toolchain quirk — a plain gcc -E on the same construct gives "passed 5 arguments, but takes just 3".

Fix (2 changes, both in strata_hip.h)

(1) cudaFuncSetAttribute back to a template function — as it was in 0.1.38's hip_compat/cuda_runtime.h:130 (which is why the wave32 backend never tripped on this):

template <typename Kernel>
inline hipError_t cudaFuncSetAttribute(Kernel kernel, hipFuncAttribute attribute, int value) {
    return hipFuncSetAttribute(reinterpret_cast<const void*>(kernel), attribute, value);
}

(2) Two definitions used beside it were missing from the platform header (present in the old one):

#define cudaFuncAttributePreferredSharedMemoryCarveout hipFuncAttributePreferredSharedMemoryCarveout
#define cudaSharedmemCarveoutMaxShared 100   // CUDA's max-carveout percentage; HIP has no constant

Scope / safety

include/strata/platform/hip_compat/ is only added to the include path inside if(STRATA_HIP_GFX906) (CMakeLists.txt:106-120), and the affected calls are inside #if defined(STRATA_HIP_GFX906). CUDA builds use the real CUDA function and the wave32 RDNA backend uses the older header — neither is touched. The change only makes the gfx906 build compile.

Verified on MI50 32 GB (IQ2_XS, 4k prompt, 256 decode, --spec 3)

build	prefill	decode
0.1.39 + this fix (clean source)	293 tok/s	42.5 tok/s
0.1.38 + local gfx906 + our patches	294 tok/s	32.4 tok/s

Correctness on both: 17*23 → 391. After the fix make is EXIT 0, the binary links libhipblas/librocblas, 108 gfx906 code objects, engine loads and serves.

Question for @JeanP00l

Thanks for the gfx906 port — it works great once building. We don't know which ROCm/toolchain you built on: on ROCm 7.2.1 the macro above is a hard error here, yet it built for you (and gcc/clang agree the construct is malformed regardless of version). Could you share your HIP/ROCm version? If your build took a path that skips these lines, that would explain it — and is worth documenting either way.

(We also have a community MI50 benchmark open as #758.)

[JeanP00l]

Thanks for the detailed report and the reproducible fix. Our gfx906 build was made with HIP/ROCm 7.14 from mixa3607/rocm-gfx906:7.14-complete, so we have not verified this build path on ROCm 7.2.1. Your explanation of the template commas being split by the function-like macro is clear, and the template wrapper is a sensible fix; the MI50 build and inference results are especially useful.

[xxDoman]

Two other open PRs fix the same two lines in this header:

• gfx906 compat: cudaFuncSetAttribute as a template function, the shared-memory carveout name (#646's fused_gr did not build) #858 — a standalone change, essentially identical to this one.
• Add experimental Linux gfx900 setup with wave64 validation #839 — the gfx900 port, which folds the same fix in among ~350 lines.

All three replace the function-like cudaFuncSetAttribute macro with a template function (the macro splits on the commas in gr_down_v3_kernel<1, 4, true>) and map cudaFuncAttributePreferredSharedMemoryCarveout. Whichever lands first, the other two will need a rebase — flagging it here so it is visible from one place. This one stays small and scoped to gfx906.

[Niko1221]

Merged for 0.1.40 (cudaFuncSetAttribute as a function on HIP/gfx906). Thank you.