ROCm build fails: host-side `rsqrtf()` in `rocm/ds4_rocm_deepseek4_vision.cuh`

[neomantra]

I've been rebasing a shared-library ds4 fork which includes CI/CD and test Docker images. My latest rebase picked up this rocm build issue that creeped up in recent work.

Posting the LLM summary, but in the end it is a one-line code change for rsqrtf.

---

Summary

make rocm (and make strix-halo) no longer compiles on ROCm 7.2.4. The
DeepSeek Vision-Exp ROCm path added in e976a54 calls rsqrtf() from host code,
but under HIP rsqrtf is only declared as a __device__ function, so hipcc
rejects the call.

Environment

• Upstream main at 6289c51 (pristine checkout, no local changes)
• Ubuntu 24.04, ROCm 7.2.4 (/opt/rocm-7.2.4, clang 22), --offload-arch=gfx1151
• make rocm ROCM_ARCH=gfx1151

Error

/opt/rocm/bin/hipcc -O3 -ffast-math -g -fno-finite-math-only -pthread -D__HIP_PLATFORM_AMD__ -Wno-unused-command-line-argument --offload-arch=gfx1151 -c -o ds4_rocm.o ds4_rocm.cu
In file included from ds4_rocm.cu:181:
./rocm/ds4_rocm_deepseek4_vision.cuh:239:25: error: no matching function for call to 'rsqrtf'
  239 |     const float alpha = rsqrtf((float)head_dim);
      |                         ^~~~~~
/opt/rocm-7.2.4/lib/llvm/lib/clang/22/include/__clang_hip_math.h:671:7: note: candidate function not viable: call to __device__ function from __host__ function
  671 | float rsqrtf(float __x) { return __ocml_rsqrt_f32(__x); }
      |       ^
1 error generated when compiling for gfx1151.
make[1]: *** [Makefile:521: ds4_rocm.o] Error 1
make: *** [Makefile:223: strix-halo] Error 2

Cause

The call is in ds4_gpu_attention_visual_mixed_batch_heads_tensor(), computing
the cuBLAS/hipBLAS alpha scale on the host before
cublasSgemmStridedBatched. CUDA's host math headers provide a host rsqrtf,
so this compiles with nvcc, but glibc has no rsqrtf and HIP only supplies the
device overload. The other rsqrtf uses in rocm/ds4_rocm_attention.cuh are
inside kernels and are fine.

Fix

Use a plain reciprocal square root for the host-side scale:

--- a/rocm/ds4_rocm_deepseek4_vision.cuh
+++ b/rocm/ds4_rocm_deepseek4_vision.cuh
@@ -236,7 +236,7 @@
     if (!cuda_ok(cudaGetLastError(), "visual attention KV pack launch"))
         return 0;
 
-    const float alpha = rsqrtf((float)head_dim);
+    const float alpha = 1.0f / sqrtf((float)head_dim);
     const float beta = 0.0f;
     cublasStatus_t status = cublasSgemmStridedBatched(

With this change make rocm ROCM_ARCH=gfx1151 builds cleanly in the same
container. Happy to send it as a PR if you prefer.