diff --git a/ggml/src/ggml-cuda/gated_delta_net.cu b/ggml/src/ggml-cuda/gated_delta_net.cu
index 5cf6968a6e1b..f0d08152acb0 100644
--- a/ggml/src/ggml-cuda/gated_delta_net.cu
+++ b/ggml/src/ggml-cuda/gated_delta_net.cu
@@ -1,6 +1,16 @@
 #include "gated_delta_net.cuh"
 #include "ggml-cuda/common.cuh"
 
+// Host and device must agree on columns/warp. Kernel uses __CUDA_ARCH__; host uses ggml_cuda_highest_compiled_arch(cc). NVIDIA Ampere+ and S_v==128 and !KDA -> 4; HIP/MUSA stay 1.
+static constexpr __host__ __device__ int gdn_cols_per_warp(int arch, int S_v, bool KDA) {
+#if defined(GGML_USE_HIP) || defined(GGML_USE_MUSA)
+    (void) arch; (void) S_v; (void) KDA;
+    return 1;
+#else
+    return GGML_CUDA_CC_IS_NVIDIA(arch) && arch >= GGML_CUDA_CC_AMPERE && S_v == 128 && !KDA ? 4 : 1;
+#endif
+}
+
 static __global__ void gdn_precompute_exp(const float * g, float * g_exp, int64_t n) {
     for (int64_t i = (int64_t) blockIdx.x*blockDim.x + threadIdx.x; i < n;
          i += (int64_t) blockDim.x*gridDim.x) {
@@ -44,10 +54,10 @@ gated_delta_net_cuda(const float * q,
     const uint32_t sequence = blockIdx.y;
     // Each warp owns one or more columns, using warp-level primitives to reduce across rows.
     const int      lane     = threadIdx.x;
-#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ == GGML_CUDA_CC_DGX_SPARK
-    constexpr int cols_per_warp = S_v == 128 && !KDA ? 4 : 1;
+#if defined(__CUDA_ARCH__)
+    constexpr int cols_per_warp = gdn_cols_per_warp(__CUDA_ARCH__, S_v, KDA);
 #else
-    constexpr int cols_per_warp = 1;
+    constexpr int cols_per_warp = 1; // host pass only; never executed
 #endif
     const int      col      = (blockIdx.z * blockDim.y + threadIdx.y) * cols_per_warp;
 
@@ -217,7 +227,8 @@ static void launch_gated_delta_net(
     const int warp_size = ggml_cuda_info().devices[ggml_cuda_get_device()].warp_size;
     const int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc;
     const int num_warps = 4;
-    const int cols_per_warp = cc == GGML_CUDA_CC_DGX_SPARK && S_v == 128 && !KDA ? 4 : 1;
+    // Same predicate as the kernel, fed the arch compiled for this cc.
+    const int cols_per_warp = GGML_CUDA_CC_IS_NVIDIA(cc) ? gdn_cols_per_warp(ggml_cuda_highest_compiled_arch(cc), S_v, KDA) : 1;
     dim3      grid_dims(H, n_seqs, (S_v + num_warps * cols_per_warp - 1) / (num_warps * cols_per_warp));
     dim3      block_dims(warp_size <= S_v ? warp_size : S_v, num_warps, 1);