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

Description

## 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`:

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

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

```cpp
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):

```cpp
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):

```cpp
#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.)*