diff --git a/rocm/ds4_rocm_matmul.cuh b/rocm/ds4_rocm_matmul.cuh
index 8a62dcd16..cdcfcbd77 100644
--- a/rocm/ds4_rocm_matmul.cuh
+++ b/rocm/ds4_rocm_matmul.cuh
@@ -275,7 +275,7 @@ static int cuda_matmul_q8_0_tensor_labeled(ds4_gpu_tensor *out, const void *mode
     }
     if (n_tok > 1) {
 #if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__)
-        if (!g_quality_mode && (in_dim % 32u) == 0u &&
+        if (g_wmma_v1 && !g_quality_mode && (in_dim % 32u) == 0u &&
             out_dim >= 1024u &&
             n_tok >= 256u &&
             in_dim <= UINT32_MAX && out_dim <= UINT32_MAX && n_tok <= UINT32_MAX) {
diff --git a/rocm/ds4_rocm_q8.cuh b/rocm/ds4_rocm_q8.cuh
index 5b2423de3..3beb06633 100644
--- a/rocm/ds4_rocm_q8.cuh
+++ b/rocm/ds4_rocm_q8.cuh
@@ -751,6 +751,9 @@ __global__ static void matmul_q8_0_f32_batch_wmma_4w_kernel(
             const _Float16 *xb = lds_x + nt * K_TILE;
             const ds4_q8_half16_t b0 = *(const ds4_q8_half16_t *)(xb);
             const ds4_q8_half16_t b1 = *(const ds4_q8_half16_t *)(xb + 16u);
+            /* gfx12xx (RDNA4) uses WMMA v2; the dispatch skips this kernel
+             * via g_wmma_v1, but the device code must compile on all arches. */
+#if !defined(__gfx1200__) && !defined(__gfx1201__)
             if (ntile == 0u) {
                 acc0 = __builtin_amdgcn_wmma_f32_16x16x16_f16_w32(a0, b0, acc0);
                 acc0 = __builtin_amdgcn_wmma_f32_16x16x16_f16_w32(a1, b1, acc0);
@@ -764,6 +767,7 @@ __global__ static void matmul_q8_0_f32_batch_wmma_4w_kernel(
                 acc3 = __builtin_amdgcn_wmma_f32_16x16x16_f16_w32(a0, b0, acc3);
                 acc3 = __builtin_amdgcn_wmma_f32_16x16x16_f16_w32(a1, b1, acc3);
             }
+#endif
         }
         __syncthreads();
     }
diff --git a/rocm/ds4_rocm_runtime.cuh b/rocm/ds4_rocm_runtime.cuh
index 07152280c..26caa7a4f 100644
--- a/rocm/ds4_rocm_runtime.cuh
+++ b/rocm/ds4_rocm_runtime.cuh
@@ -21,6 +21,9 @@ static int g_cublas_ready;
 #include "ds4_rocm_hipblaslt.cuh"
 #endif
 static int g_quality_mode;
+/* 1 if device supports RDNA3 WMMA v1 (__builtin_amdgcn_wmma_*_w32).
+ * 0 for RDNA4 (gfx12xx) which uses WMMA v2 — skip the _4w kernel. */
+static int g_wmma_v1;
 
 enum {
     DS4_ROCM_N_EXPERT = 256u,
@@ -5725,6 +5728,15 @@ extern "C" int ds4_gpu_init(void) {
     if (cudaGetDeviceProperties(&prop, dev) == cudaSuccess) {
         fprintf(stderr, DS4_GPU_LOG_PREFIX "backend initialized on %s (sm_%d%d)\n",
                 prop.name, prop.major, prop.minor);
+#ifdef __HIP_PLATFORM_AMD__
+        /* RDNA4 (gfx12xx) needs WMMA v2; gfx11xx (RDNA3/3.5) uses v1 builtins. */
+        g_wmma_v1 = (strstr(prop.gcnArchName, "gfx12") == NULL);
+        fprintf(stderr, DS4_GPU_LOG_PREFIX "WMMA v1 support: %s (arch: %s)\n",
+                g_wmma_v1 ? "yes" : "no (gfx12 RDNA4, using non-WMMA path)",
+                prop.gcnArchName);
+#else
+        g_wmma_v1 = 1;
+#endif
     }
     if (!g_cublas_ready) {
         if (!cublas_ok(cublasCreate(&g_cublas), "create handle")) return 0;