fix(ggml-cuda): restore fork build portability across CUDA + HIP backends

The fork's ggml-cuda failed to build on the CUDA host (nvcc, GGML_CUDA=ON)
and on gfx803/HIP — AMD-only vocabulary and a missing warp-sync shim had
leaked into shared/core files. Backend-portability only; no quant math changed.

- set-rows.cu: fp8 e4m3 device decode used __builtin_amdgcn_ldexp
  unconditionally (undefined under nvcc). Guard #if GGML_USE_HIP (amdgcn) /
  #else portable ldexpf — identical numerics, AMD codegen unchanged.

- turbo_fp8_hadamard.cuh: mt_turbo_fp8_fwht{,_half} declared with raw
  hipError_t/hipStream_t, tripping CUDA_CHECK (cudaError_t) at ml8.cu:1533.
  Switch to CUDA vocabulary; vendors/hip.h remaps it back to hip* under HIP.

- vendors/hip.h:
  * add cudaErrorInvalidValue -> hipErrorInvalidValue remap (needed above).
  * add __syncwarp(...) -> __builtin_amdgcn_wave_barrier() shim, beside the
    existing __shfl_*_sync shims. __syncwarp is undeclared in the HIP
    toolchain for TUs that don't pull it in (e.g. mt_pagedattn_decode.cu);
    the shim makes that warp barrier portable WITHOUT editing that file, and
    PRESERVES the call on every arch incl. RDNA4 (gfx1201).

Verified:
- CUDA sm_61 (GTX 1070): full llama-server links + runs on-GPU (token output).
- HIP gfx803 / ROCm 6.4 (RX 480): full llama-server links + runs on-GPU
  (turbo4 KV, 6.18 GB VRAM, token output).
- mt_pagedattn_decode.cu + fattn-vec compile clean for gfx803 AND gfx1201
  (-fsyntax-only); RDNA4 WMMA decode path unchanged.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

Author: kmbandy

ggml/src/ggml-cuda/set-rows.cu

@@ -1167,7 +1167,11 @@ static __device__ __forceinline__ float fp8_e4m3_byte_to_fp32_dev(uint8_t b) {
     int e    = (b >> 3) & 0xF;
     int m    = b & 0x7;
     float v  = (e == 0) ? (1.0f / 64.0f) * (m / 8.0f)
+#if defined(GGML_USE_HIP)
                         : __builtin_amdgcn_ldexp(1.0f + m / 8.0f, e - 7);
+#else
+                        : ldexpf(1.0f + m / 8.0f, e - 7);
+#endif
     return sign ? -v : v;
 }
 

ggml/src/ggml-cuda/turbo_fp8_hadamard.cuh

@@ -95,14 +95,14 @@ __global__ void mt_turbo_fp8_fwht_kernel(float * __restrict__ data, int n_rows,
 //
 // row_stride is in fp32 elements (typically == head_dim for tightly packed rows).
 // ─────────────────────────────────────────────────────────────────────────
-static inline hipError_t mt_turbo_fp8_fwht(
-    hipStream_t stream,
+static inline cudaError_t mt_turbo_fp8_fwht(
+    cudaStream_t stream,
     float *     data_device,
     int         n_rows,
     int         head_dim,
     int         row_stride
 ) {
-    if (n_rows <= 0) return hipSuccess;
+    if (n_rows <= 0) return cudaSuccess;
     dim3 grid(n_rows, 1, 1);
 
     switch (head_dim) {
@@ -128,9 +128,9 @@ static inline hipError_t mt_turbo_fp8_fwht(
             mt_turbo_fp8_fwht_kernel<1024><<<grid, dim3(1024),0, stream>>>(data_device, n_rows, row_stride);
             break;
         default:
-            return hipErrorInvalidValue;
+            return cudaErrorInvalidValue;
     }
-    return hipGetLastError();
+    return cudaGetLastError();
 }
 
 
@@ -176,14 +176,14 @@ __global__ void mt_turbo_fp8_fwht_half_kernel(__half * __restrict__ data, int n_
     row_ptr[tid] = __float2half(smem[tid] * scale);
 }
 
-static inline hipError_t mt_turbo_fp8_fwht_half(
-    hipStream_t stream,
+static inline cudaError_t mt_turbo_fp8_fwht_half(
+    cudaStream_t stream,
     __half *    data_device,
     int         n_rows,
     int         head_dim,
     int         row_stride
 ) {
-    if (n_rows <= 0) return hipSuccess;
+    if (n_rows <= 0) return cudaSuccess;
     dim3 grid(n_rows, 1, 1);
 
     switch (head_dim) {
@@ -194,9 +194,9 @@ static inline hipError_t mt_turbo_fp8_fwht_half(
         case 256:  mt_turbo_fp8_fwht_half_kernel<256> <<<grid, dim3(256), 0, stream>>>(data_device, n_rows, row_stride); break;
         case 512:  mt_turbo_fp8_fwht_half_kernel<512> <<<grid, dim3(512), 0, stream>>>(data_device, n_rows, row_stride); break;
         case 1024: mt_turbo_fp8_fwht_half_kernel<1024><<<grid, dim3(1024),0, stream>>>(data_device, n_rows, row_stride); break;
-        default: return hipErrorInvalidValue;
+        default: return cudaErrorInvalidValue;
     }
-    return hipGetLastError();
+    return cudaGetLastError();
 }
 
 

ggml/src/ggml-cuda/vendors/hip.h

@@ -52,6 +52,11 @@
 #define __SHFL_XOR_GET(_1, _2, _3, _4, NAME, ...) NAME
 #define __shfl_xor_sync(...) __SHFL_XOR_GET(__VA_ARGS__, __SHFL_XOR_SYNC_4, __SHFL_XOR_SYNC_3)(__VA_ARGS__)
 
+// __syncwarp: HIP/ROCm has no warp-mask sync intrinsic (undeclared on gfx803 and
+// in this toolchain on RDNA arches). A wavefront execution barrier is the
+// equivalent; HIP ignores the lane mask, like the __shfl_*_sync shims above.
+#define __syncwarp(...) __builtin_amdgcn_wave_barrier()
+
 // __shfl_down_sync: support 3-arg and 4-arg calls (HIP ignores mask)
 #define __SHFL_DOWN_SYNC_3(mask, var, delta)        __shfl_down(var, delta, warpSize)
 #define __SHFL_DOWN_SYNC_4(mask, var, delta, width) __shfl_down(var, delta, width)
@@ -82,6 +87,7 @@
 #define cudaDeviceProp hipDeviceProp_t
 #define cudaDeviceSynchronize hipDeviceSynchronize
 #define cudaError_t hipError_t
+#define cudaErrorInvalidValue hipErrorInvalidValue
 #define cudaErrorMemoryAllocation hipErrorOutOfMemory
 #define cudaErrorPeerAccessAlreadyEnabled hipErrorPeerAccessAlreadyEnabled
 #define cudaErrorPeerAccessNotEnabled hipErrorPeerAccessNotEnabled