Skip to content

rocm: fix build on gfx1201 (RDNA4, R9700) — gate the WMMA v1 kernel #599

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
cm999club wants to merge 1 commit into
base: main
Choose a base branch
from
+17 −1
Open

rocm: fix build on gfx1201 (RDNA4, R9700) — gate the WMMA v1 kernel #599

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion rocm/ds4_rocm_matmul.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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) {
Expand Down
4 changes: 4 additions & 0 deletions rocm/ds4_rocm_q8.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand All @@ -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();
}
Expand Down
12 changes: 12 additions & 0 deletions rocm/ds4_rocm_runtime.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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,
Expand Down Expand Up @@ -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;
Expand Down