Skip to content

Commit ef90ac0

Browse files
superm1superm1
committed
ggml: add RDNA4m support for gfx117x
gfx117x is more like RDNA4, with 128b WMMA. But they're not identical. Adjust the defines accordingly to have a new RDNA4m and use it where it makes sense. Assisted-by: Codex
1 parent 1877c1b commit ef90ac0

9 files changed

Lines changed: 71 additions & 52 deletions

File tree

‎ggml/src/ggml-cuda/common.cuh‎

Lines changed: 18 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -81,15 +81,17 @@
8181
#define GGML_CUDA_CC_RDNA2 (GGML_CUDA_CC_OFFSET_AMD + 0x1030) // RX 6000, minimum for dp4a
8282
#define GGML_CUDA_CC_RDNA3 (GGML_CUDA_CC_OFFSET_AMD + 0x1100) // RX 7000, minimum for WMMA
8383
#define GGML_CUDA_CC_RDNA3_5 (GGML_CUDA_CC_OFFSET_AMD + 0x1150) // AI 370, AI Max 395 laptops.
84+
#define GGML_CUDA_CC_RDNA4_M (GGML_CUDA_CC_OFFSET_AMD + 0x1170) // gfx117x
8485
#define GGML_CUDA_CC_RDNA4 (GGML_CUDA_CC_OFFSET_AMD + 0x1200) // RX 9000
8586

8687
#define GGML_CUDA_CC_IS_AMD(cc) (cc >= GGML_CUDA_CC_OFFSET_AMD)
8788
#define GGML_CUDA_CC_IS_RDNA(cc) (cc >= GGML_CUDA_CC_RDNA1)
8889
#define GGML_CUDA_CC_IS_RDNA1(cc) (cc >= GGML_CUDA_CC_RDNA1 && cc < GGML_CUDA_CC_RDNA2)
8990
#define GGML_CUDA_CC_IS_RDNA2(cc) (cc >= GGML_CUDA_CC_RDNA2 && cc < GGML_CUDA_CC_RDNA3)
9091
#define GGML_CUDA_CC_IS_RDNA3_0(cc) (cc >= GGML_CUDA_CC_RDNA3 && cc < GGML_CUDA_CC_RDNA3_5)
91-
#define GGML_CUDA_CC_IS_RDNA3_5(cc) (cc >= GGML_CUDA_CC_RDNA3_5 && cc < GGML_CUDA_CC_RDNA4)
92+
#define GGML_CUDA_CC_IS_RDNA3_5(cc) (cc >= GGML_CUDA_CC_RDNA3_5 && cc < GGML_CUDA_CC_RDNA4_M)
9293
#define GGML_CUDA_CC_IS_RDNA3(cc) (GGML_CUDA_CC_IS_RDNA3_0(cc) || GGML_CUDA_CC_IS_RDNA3_5(cc))
94+
#define GGML_CUDA_CC_IS_RDNA4_M(cc) (cc >= GGML_CUDA_CC_RDNA4_M && cc < GGML_CUDA_CC_RDNA4)
9395
#define GGML_CUDA_CC_IS_RDNA4(cc) (cc >= GGML_CUDA_CC_RDNA4)
9496
#define GGML_CUDA_CC_IS_GCN_APU(cc) ((cc) == GGML_CUDA_CC_GFX909 || (cc) == GGML_CUDA_CC_GFX90C)
9597
#define GGML_CUDA_CC_IS_GCN(cc) ((cc > GGML_CUDA_CC_OFFSET_AMD && cc < GGML_CUDA_CC_CDNA1) || GGML_CUDA_CC_IS_GCN_APU(cc))
@@ -276,9 +278,17 @@ static const char * cu_get_error_str(CUresult err) {
276278
#define AMD_MFMA_AVAILABLE
277279
#endif // defined(GGML_USE_HIP) && defined(CDNA) && !defined(GGML_HIP_NO_MMQ_MFMA)
278280

279-
#if defined(GGML_USE_HIP) && (defined(RDNA4) || defined(RDNA3))
281+
#if defined(GGML_USE_HIP) && (defined(RDNA4) || defined(RDNA4_M))
282+
#define AMD_WMMA_128B_AVAILABLE
283+
#endif // defined(GGML_USE_HIP) && (defined(RDNA4) || defined(RDNA4_M))
284+
285+
#if defined(GGML_USE_HIP) && defined(RDNA3)
286+
#define AMD_WMMA_256B_AVAILABLE
287+
#endif // defined(GGML_USE_HIP) && defined(RDNA3)
288+
289+
#if defined(AMD_WMMA_128B_AVAILABLE) || defined(AMD_WMMA_256B_AVAILABLE)
280290
#define AMD_WMMA_AVAILABLE
281-
#endif // defined(GGML_USE_HIP) && defined(RDNA4)
291+
#endif // defined(AMD_WMMA_128B_AVAILABLE) || defined(AMD_WMMA_256B_AVAILABLE)
282292

283293
// The Volta instructions are in principle available on Turing or newer but they are effectively unusable:
284294
#if !defined(GGML_USE_HIP) && __CUDA_ARCH__ == GGML_CUDA_CC_VOLTA
@@ -325,7 +335,7 @@ static bool fast_fp16_hardware_available(const int cc) {
325335
// To be used for feature selection of external libraries, e.g. cuBLAS.
326336
static bool fp16_mma_hardware_available(const int cc) {
327337
return (GGML_CUDA_CC_IS_NVIDIA(cc) && cc >= GGML_CUDA_CC_VOLTA) ||
328-
GGML_CUDA_CC_IS_CDNA(cc) || GGML_CUDA_CC_IS_RDNA3(cc) || GGML_CUDA_CC_IS_RDNA4(cc) ||
338+
GGML_CUDA_CC_IS_CDNA(cc) || GGML_CUDA_CC_IS_RDNA3(cc) || GGML_CUDA_CC_IS_RDNA4_M(cc) || GGML_CUDA_CC_IS_RDNA4(cc) ||
329339
(GGML_CUDA_CC_IS_MTHREADS(cc) && cc >= GGML_CUDA_CC_QY2);
330340
}
331341

@@ -348,7 +358,7 @@ static bool amd_mfma_available(const int cc) {
348358
}
349359

350360
static bool amd_wmma_available(const int cc) {
351-
return (GGML_CUDA_CC_IS_RDNA4(cc) || GGML_CUDA_CC_IS_RDNA3(cc));
361+
return GGML_CUDA_CC_IS_RDNA3(cc) || GGML_CUDA_CC_IS_RDNA4_M(cc) || GGML_CUDA_CC_IS_RDNA4(cc);
352362
}
353363

354364
static bool volta_mma_available(const int cc) {
@@ -715,7 +725,7 @@ static __device__ __forceinline__ int ggml_cuda_dp4a(const int a, const int b, i
715725
#if defined(GGML_USE_HIP)
716726
#if defined(CDNA) || defined(RDNA2) || defined(__gfx906__)
717727
c = __builtin_amdgcn_sdot4(a, b, c, false);
718-
#elif defined(RDNA3) || defined(RDNA4)
728+
#elif defined(RDNA3) || defined(RDNA4_M) || defined(RDNA4)
719729
c = __builtin_amdgcn_sudot4( true, a, true, b, c, false);
720730
#elif defined(RDNA1) || defined(__gfx900__)
721731
int tmp1;
@@ -760,9 +770,9 @@ static __device__ __forceinline__ void ggml_cuda_mad(float & acc, const float2 v
760770
acc += v.y*u.y;
761771
}
762772

763-
#if defined(GGML_USE_HIP) && (defined(RDNA2) || defined(RDNA3) || defined(RDNA4) || defined(__gfx906__) || defined(CDNA))
773+
#if defined(GGML_USE_HIP) && (defined(RDNA2) || defined(RDNA3) || defined(RDNA4_M) || defined(RDNA4) || defined(__gfx906__) || defined(CDNA))
764774
#define V_DOT2_F32_F16_AVAILABLE
765-
#endif // defined(GGML_USE_HIP) && (defined(RDNA2) || defined(RDNA3) || defined(RDNA4) || defined(__gfx906__) || defined(CDNA))
775+
#endif // defined(GGML_USE_HIP) && (defined(RDNA2) || defined(RDNA3) || defined(RDNA4_M) || defined(RDNA4) || defined(__gfx906__) || defined(CDNA))
766776

767777
static __device__ __forceinline__ void ggml_cuda_mad(float & acc, const half2 v, const half2 u) {
768778
#ifdef V_DOT2_F32_F16_AVAILABLE

‎ggml/src/ggml-cuda/fattn-mma-f16.cuh‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1225,7 +1225,7 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile(
12251225
T_B_KQ Q_B[(Q_in_reg ? DKQ/(2*T_B_KQ::J) : 1)];
12261226
#if defined(TURING_MMA_AVAILABLE)
12271227
T_C_VKQ VKQ_C[cols_per_warp == 8 ? DV/T_C_VKQ::I : DV/(2*T_C_VKQ::J)];
1228-
#elif defined(AMD_WMMA_AVAILABLE) && defined(RDNA3)
1228+
#elif defined(AMD_WMMA_256B_AVAILABLE)
12291229
T_C_VKQ VKQ_C[DV % 32 != 0 ? DV/T_C_VKQ::J : DV/(2*T_C_VKQ::J)];
12301230
#elif defined(AMD_WMMA_AVAILABLE) || defined(AMD_MFMA_AVAILABLE)
12311231
T_C_VKQ VKQ_C[ DV/(2*T_C_VKQ::J)];

‎ggml/src/ggml-cuda/ggml-cuda.cu‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1511,7 +1511,7 @@ static void ggml_cuda_mul_mat_cublas_impl(ggml_backend_cuda_context & ctx, const
15111511
if (compute_type == GGML_TYPE_F16) {
15121512
prefer_f32_output = cc == GGML_CUDA_CC_VOLTA || GGML_CUDA_CC_IS_RDNA4(cc) || GGML_CUDA_CC_IS_CDNA(cc);
15131513
} else if (compute_type == GGML_TYPE_BF16) {
1514-
prefer_f32_output = !GGML_CUDA_CC_IS_RDNA3(cc) && !GGML_CUDA_CC_IS_CDNA(cc);
1514+
prefer_f32_output = !GGML_CUDA_CC_IS_RDNA3(cc) && !GGML_CUDA_CC_IS_RDNA4_M(cc) && !GGML_CUDA_CC_IS_CDNA(cc);
15151515
}
15161516

15171517
if (prefer_f32_output) {

‎ggml/src/ggml-cuda/mma.cuh‎

Lines changed: 31 additions & 31 deletions
Original file line numberDiff line numberDiff line change
@@ -76,23 +76,23 @@ namespace ggml_cuda_mma {
7676
// For the A/C matrices this means I major == row major, J major == column major.
7777
// For the B matrix this means I major == column major, J major == row major.
7878
// MIRRORED == Each data value is held exactly once per thread subgroup.
79-
DATA_LAYOUT_I_MAJOR = 0, // Always used for Turing, Ampere, Ada Lovelace, consumer Blackwell, matrix A&B for RDNA4 and CDNA.
80-
DATA_LAYOUT_J_MAJOR = 10, // Matrix C for CDNA and RDNA4, int and float matrix C for RDNA3.
79+
DATA_LAYOUT_I_MAJOR = 0, // Always used for Turing, Ampere, Ada Lovelace, consumer Blackwell, matrix A&B for WMMA128 and CDNA.
80+
DATA_LAYOUT_J_MAJOR = 10, // Matrix C for CDNA and WMMA128, int and float matrix C for WMMA256.
8181
DATA_LAYOUT_I_MAJOR_MIRRORED = 20, // Volta, matrix A&B for RDNA3.
8282
DATA_LAYOUT_J_MAJOR_MIRRORED = 30,
83-
DATA_LAYOUT_I_MAJOR_SCRAMBLED = 40, // Scrambled matrix C for faster transposition (RDNA4/CDNA), convert to float to unscramble.
83+
DATA_LAYOUT_I_MAJOR_SCRAMBLED = 40, // Scrambled matrix C for faster transposition (WMMA128/CDNA), convert to float to unscramble.
8484
};
8585
// Implemented mma combinations are:
8686
// - (I_MAJOR, I_MAJOR) -> I_MAJOR
8787
// - (I_MAJOR, I_MAJOR_MIRRORED) -> I_MAJOR
8888
// - (I_MAJOR, J_MAJOR_MIRRORED) -> I_MAJOR
8989

9090
static constexpr __device__ data_layout get_input_data_layout() {
91-
#if defined(RDNA3) || defined(VOLTA_MMA_AVAILABLE)
91+
#if defined(AMD_WMMA_256B_AVAILABLE) || defined(VOLTA_MMA_AVAILABLE)
9292
return DATA_LAYOUT_I_MAJOR_MIRRORED;
9393
#else
9494
return DATA_LAYOUT_I_MAJOR;
95-
#endif // defined(RDNA3) || defined(VOLTA_MMA_AVAILABLE)
95+
#endif // defined(AMD_WMMA_256B_AVAILABLE) || defined(VOLTA_MMA_AVAILABLE)
9696
}
9797

9898
template <int I_, int J_, typename T, data_layout ds_=DATA_LAYOUT_I_MAJOR>
@@ -201,7 +201,7 @@ namespace ggml_cuda_mma {
201201

202202
static __device__ __forceinline__ int get_j(const int l) {
203203
if constexpr (I == 16 && J == 16) {
204-
#if defined(RDNA3)
204+
#if defined(AMD_WMMA_256B_AVAILABLE)
205205
if constexpr (std::is_same_v<T, float> || std::is_same_v<T, int>) {
206206
// matrix C
207207
return 2 * l + (threadIdx.x / 16);
@@ -210,11 +210,11 @@ namespace ggml_cuda_mma {
210210
return l;
211211
}
212212
#else
213-
// matrix C is the transposed matrix A&B on RDNA4
213+
// matrix C is the transposed matrix A&B on WMMA128
214214
return ne * (threadIdx.x / 16) + l;
215-
#endif // defined(RDNA3)
215+
#endif // defined(AMD_WMMA_256B_AVAILABLE)
216216
} else if constexpr (I == 16 && J == 8) {
217-
// mmq input for RDNA4
217+
// mmq input for WMMA128
218218
return ne * (threadIdx.x / 16) + l;
219219
} else if constexpr (I == 16 && J == 4) {
220220
return ne * (threadIdx.x / 16) + l;
@@ -573,7 +573,7 @@ namespace ggml_cuda_mma {
573573
static constexpr int I = I_;
574574
static constexpr int J = J_;
575575
static constexpr data_layout dl = DATA_LAYOUT_I_MAJOR_MIRRORED;
576-
#if defined(RDNA3)
576+
#if defined(AMD_WMMA_256B_AVAILABLE)
577577
static constexpr int ne = tile<I_, J_, float, DATA_LAYOUT_I_MAJOR_MIRRORED>::ne;
578578

579579
half2 x[ne] = {{0.0f, 0.0f}};
@@ -616,7 +616,7 @@ namespace ggml_cuda_mma {
616616
return -1;
617617
}
618618
}
619-
#endif // defined(RDNA3)
619+
#endif // defined(AMD_WMMA_256B_AVAILABLE)
620620
};
621621

622622
template <int I_, int J_>
@@ -694,7 +694,7 @@ namespace ggml_cuda_mma {
694694
};
695695

696696
static __device__ __forceinline__ tile<16, 16, half2, DATA_LAYOUT_I_MAJOR> unscramble(const tile<16, 16, half2, DATA_LAYOUT_I_MAJOR_SCRAMBLED> & t) {
697-
#if defined(AMD_MFMA_AVAILABLE) || (defined(AMD_WMMA_AVAILABLE) && defined(RDNA4))
697+
#if defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_128B_AVAILABLE)
698698
tile<16, 16, half2, DATA_LAYOUT_I_MAJOR> ret;
699699
#pragma unroll
700700
for (int l0 = 0; l0 < t.ne/2; ++l0) {
@@ -705,7 +705,7 @@ namespace ggml_cuda_mma {
705705
#else
706706
NO_DEVICE_CODE;
707707
GGML_UNUSED(t);
708-
#endif // defined(AMD_MFMA_AVAILABLE) || (defined(AMD_WMMA_AVAILABLE) && defined(RDNA4))
708+
#endif // defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_128B_AVAILABLE)
709709
}
710710

711711
#if defined(TURING_MMA_AVAILABLE)
@@ -726,7 +726,7 @@ namespace ggml_cuda_mma {
726726

727727
return ret;
728728
}
729-
#elif defined(AMD_WMMA_AVAILABLE) && defined(RDNA3)
729+
#elif defined(AMD_WMMA_256B_AVAILABLE)
730730
static __device__ __forceinline__ tile<16, 8, half2, DATA_LAYOUT_I_MAJOR_MIRRORED> get_half2(
731731
const tile<16, 16, float, DATA_LAYOUT_I_MAJOR> & tile_float) {
732732
tile<16, 8, half2, DATA_LAYOUT_I_MAJOR_MIRRORED> ret;
@@ -1021,7 +1021,7 @@ namespace ggml_cuda_mma {
10211021
: "r"(Axi[2]), "r"(Axi[3]), "r"(Bxi[3]));
10221022
#endif // __CUDA_ARCH__ >= GGML_CUDA_CC_AMPERE
10231023
#elif defined(AMD_WMMA_AVAILABLE)
1024-
#if defined(RDNA4)
1024+
#if defined(AMD_WMMA_128B_AVAILABLE)
10251025
using halfx8_t = __attribute__((ext_vector_type(8))) _Float16;
10261026
halfx8_t& acc_frag = reinterpret_cast<halfx8_t&>(D.x[0]);
10271027
const halfx8_t& a_frag = reinterpret_cast<const halfx8_t&>(A.x[0]);
@@ -1030,7 +1030,7 @@ namespace ggml_cuda_mma {
10301030
#else
10311031
GGML_UNUSED_VARS(D, A, B);
10321032
NO_DEVICE_CODE;
1033-
#endif // defined(RDNA4)
1033+
#endif // defined(AMD_WMMA_128B_AVAILABLE)
10341034
#elif defined(AMD_MFMA_AVAILABLE)
10351035
// MFMA: FP16 input, FP32 accumulate, convert back to half2.
10361036
using halfx4_t = __attribute__((ext_vector_type(4))) _Float16;
@@ -1068,15 +1068,15 @@ namespace ggml_cuda_mma {
10681068
static __device__ __forceinline__ void mma(
10691069
tile<16, 16, half2, DATA_LAYOUT_I_MAJOR_SCRAMBLED> & D, const tile<32, 8, half2, DATA_LAYOUT_I_MAJOR> & A,
10701070
const tile<16, 8, half2, DATA_LAYOUT_I_MAJOR> & B) {
1071-
#if defined(AMD_MFMA_AVAILABLE) || (defined(AMD_WMMA_AVAILABLE) && defined(RDNA4))
1071+
#if defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_128B_AVAILABLE)
10721072
tile<16, 8, half2> * D16 = (tile<16, 8, half2> *) &D;
10731073
const tile<16, 8, half2> * A16 = (const tile<16, 8, half2> *) &A;
10741074
mma(D16[0], A16[0], B);
10751075
mma(D16[1], A16[1], B);
10761076
#else
10771077
GGML_UNUSED_VARS(D, A, B);
10781078
NO_DEVICE_CODE;
1079-
#endif // defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) && defined(RDNA4)
1079+
#endif // defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_128B_AVAILABLE)
10801080
}
10811081

10821082
template <data_layout dl_ab, data_layout dl_d>
@@ -1223,14 +1223,14 @@ namespace ggml_cuda_mma {
12231223
: "r"(Axi[2]), "r"(Axi[3]), "r"(Bxi[3]));
12241224
#endif // __CUDA_ARCH__ >= GGML_CUDA_CC_AMPERE
12251225
#elif defined(AMD_WMMA_AVAILABLE)
1226-
#if defined(RDNA4)
1226+
#if defined(AMD_WMMA_128B_AVAILABLE)
12271227
using halfx8_t = __attribute__((ext_vector_type(8))) _Float16;
12281228
using floatx8_t = __attribute__((ext_vector_type(8))) float;
12291229
floatx8_t& acc_frag = reinterpret_cast<floatx8_t&>(D.x[0]);
12301230
const halfx8_t& a_frag = reinterpret_cast<const halfx8_t&>(A.x[0]);
12311231
const halfx8_t& b_frag = reinterpret_cast<const halfx8_t&>(B.x[0]);
12321232
acc_frag = __builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12(a_frag, b_frag, acc_frag);
1233-
#elif defined(RDNA3)
1233+
#elif defined(AMD_WMMA_256B_AVAILABLE)
12341234
using halfx16_t = __attribute__((ext_vector_type(16))) _Float16;
12351235
using floatx8_t = __attribute__((ext_vector_type(8))) float;
12361236
floatx8_t& acc_frag = reinterpret_cast<floatx8_t&>(D.x[0]);
@@ -1240,7 +1240,7 @@ namespace ggml_cuda_mma {
12401240
#else
12411241
GGML_UNUSED_VARS(D, A, B);
12421242
NO_DEVICE_CODE;
1243-
#endif // RDNA4
1243+
#endif // AMD_WMMA_128B_AVAILABLE
12441244
#elif defined(AMD_MFMA_AVAILABLE)
12451245
using halfx4_t = __attribute__((ext_vector_type(4))) _Float16;
12461246
using floatx4_t = __attribute__((ext_vector_type(4))) float;
@@ -1258,14 +1258,14 @@ namespace ggml_cuda_mma {
12581258
static __device__ __forceinline__ void mma(
12591259
tile<16, 16, float, dl_d> & D, const tile<16, 8, nv_bfloat162, dl_ab> & A, const tile<16, 8, nv_bfloat162, dl_ab> & B) {
12601260
#if defined(AMD_WMMA_AVAILABLE)
1261-
#if defined(RDNA4)
1261+
#if defined(AMD_WMMA_128B_AVAILABLE)
12621262
using bf16x8_t = __attribute__((ext_vector_type(8))) __bf16;
12631263
using floatx8_t = __attribute__((ext_vector_type(8))) float;
12641264
floatx8_t& acc_frag = reinterpret_cast<floatx8_t&>(D.x[0]);
12651265
const bf16x8_t& a_frag = reinterpret_cast<const bf16x8_t&>(A.x[0]);
12661266
const bf16x8_t& b_frag = reinterpret_cast<const bf16x8_t&>(B.x[0]);
12671267
acc_frag = __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32_gfx12(a_frag, b_frag, acc_frag);
1268-
#elif defined(RDNA3)
1268+
#elif defined(AMD_WMMA_256B_AVAILABLE)
12691269
using bf16x16_t = __attribute__((ext_vector_type(16))) __bf16;
12701270
using floatx8_t = __attribute__((ext_vector_type(8))) float;
12711271
floatx8_t& acc_frag = reinterpret_cast<floatx8_t&>(D.x[0]);
@@ -1275,7 +1275,7 @@ namespace ggml_cuda_mma {
12751275
#else
12761276
GGML_UNUSED_VARS(D, A, B);
12771277
NO_DEVICE_CODE;
1278-
#endif // defined(RDNA4)
1278+
#endif // defined(AMD_WMMA_128B_AVAILABLE)
12791279
#elif defined(AMD_MFMA_AVAILABLE)
12801280
using floatx4_t = __attribute__((ext_vector_type(4))) float;
12811281
floatx4_t& acc_frag = reinterpret_cast<floatx4_t&>(D.x[0]);
@@ -1317,19 +1317,19 @@ namespace ggml_cuda_mma {
13171317
#elif defined(AMD_WMMA_AVAILABLE)
13181318
using int32x8_t = __attribute__((__vector_size__(8 * sizeof(int)))) int;
13191319
int32x8_t * acc = (int32x8_t *) D.x;
1320-
#if defined(RDNA4)
1320+
#if defined(AMD_WMMA_128B_AVAILABLE)
13211321
using int32x2_t = __attribute__((__vector_size__(2 * sizeof(int)))) int;
13221322
int32x2_t * a_vec = (int32x2_t *) A.x;
13231323
int32x2_t * b_vec = (int32x2_t *) B.x;
13241324
acc[0] = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12(true, a_vec[0], true, b_vec[0], acc[0], true);
13251325
acc[0] = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12(true, a_vec[1], true, b_vec[1], acc[0], true);
1326-
#elif defined(RDNA3)
1326+
#elif defined(AMD_WMMA_256B_AVAILABLE)
13271327
using int32x4_t = __attribute__((__vector_size__(4 * sizeof(int)))) int;
13281328
int32x4_t * a_vec = (int32x4_t *) A.x;
13291329
int32x4_t * b_vec = (int32x4_t *) B.x;
13301330
acc[0] = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32(true, a_vec[0], true, b_vec[0], acc[0], true);
13311331
acc[0] = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32(true, a_vec[1], true, b_vec[1], acc[0], true);
1332-
#endif // RDNA4
1332+
#endif // AMD_WMMA_128B_AVAILABLE
13331333
#else
13341334
GGML_UNUSED_VARS(D, A, B);
13351335
NO_DEVICE_CODE;
@@ -1406,7 +1406,7 @@ namespace ggml_cuda_mma {
14061406
static __device__ __forceinline__ void mma(
14071407
tile<16, 16, half2, DATA_LAYOUT_I_MAJOR> & D, const tile<32, 8, half2, DATA_LAYOUT_I_MAJOR_MIRRORED> & A,
14081408
const tile<16, 8, half2, DATA_LAYOUT_I_MAJOR_MIRRORED> & B) {
1409-
#if defined(AMD_WMMA_AVAILABLE) && defined(RDNA3)
1409+
#if defined(AMD_WMMA_256B_AVAILABLE)
14101410
using halfx16_t = __attribute__((ext_vector_type(16))) _Float16;
14111411
halfx16_t * xD = (halfx16_t *) D.x;
14121412
const halfx16_t * xA = (const halfx16_t *) A.x;
@@ -1435,17 +1435,17 @@ namespace ggml_cuda_mma {
14351435
#elif defined(AMD_WMMA_AVAILABLE)
14361436
using int32x8_t = __attribute__((__vector_size__(8 * sizeof(int)))) int;
14371437
int32x8_t * acc = (int32x8_t *) D.x;
1438-
#if defined(RDNA4)
1438+
#if defined(AMD_WMMA_128B_AVAILABLE)
14391439
using int32x2_t = __attribute__((__vector_size__(2 * sizeof(int)))) int;
14401440
int32x2_t * a_vec = (int32x2_t *) A.x;
14411441
int32x2_t * b_vec = (int32x2_t *) B.x;
14421442
acc[0] = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12(true, a_vec[0], true, b_vec[0], acc[0], false);
1443-
#elif defined(RDNA3)
1443+
#elif defined(AMD_WMMA_256B_AVAILABLE)
14441444
using int32x4_t = __attribute__((__vector_size__(4 * sizeof(int)))) int;
14451445
int32x4_t * a_vec = (int32x4_t *) A.x;
14461446
int32x4_t * b_vec = (int32x4_t *) B.x;
14471447
acc[0] = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32(true, a_vec[0], true, b_vec[0], acc[0], false);
1448-
#endif // RDNA4
1448+
#endif // AMD_WMMA_128B_AVAILABLE
14491449
#else
14501450
GGML_UNUSED(D);
14511451
GGML_UNUSED(A);

‎ggml/src/ggml-cuda/mmq.cu‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -354,7 +354,7 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t
354354
}
355355

356356
if (amd_wmma_available(cc)) {
357-
if (GGML_CUDA_CC_IS_RDNA3(cc)) {
357+
if (GGML_CUDA_CC_IS_RDNA3(cc) || GGML_CUDA_CC_IS_RDNA4_M(cc)) {
358358
// High expert counts are almost always better on MMQ due to
359359
// the synchronization overhead in the cuBLAS/hipBLAS path:
360360
// https://github.com/ggml-org/llama.cpp/pull/18202
@@ -371,7 +371,7 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t
371371
return ne11 <= (GGML_CUDA_CC_IS_RDNA3_0(cc) ? 128 : 256);
372372
case GGML_TYPE_IQ2_XS:
373373
case GGML_TYPE_IQ2_S:
374-
return GGML_CUDA_CC_IS_RDNA3_5(cc) || ne11 <= 128;
374+
return GGML_CUDA_CC_IS_RDNA3_5(cc) || GGML_CUDA_CC_IS_RDNA4_M(cc) || ne11 <= 128;
375375
default:
376376
return true;
377377
}

‎ggml/src/ggml-cuda/mmq.cuh‎

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -234,6 +234,9 @@ static __host__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config(const ggml_type ty
234234
if (GGML_CUDA_CC_IS_RDNA4(cc)) {
235235
return ggml_cuda_mmq_get_config_rdna4(type, J, fallback);
236236
}
237+
if (GGML_CUDA_CC_IS_RDNA4_M(cc)) {
238+
return ggml_cuda_mmq_get_config_rdna3_5(type, J, fallback);
239+
}
237240
if (GGML_CUDA_CC_IS_RDNA3_5(cc)) {
238241
return ggml_cuda_mmq_get_config_rdna3_5(type, J, fallback);
239242
}
@@ -260,6 +263,8 @@ static constexpr __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config(ggml_t
260263
return ggml_cuda_mmq_get_config_cdna(type, J, fallback);
261264
#elif defined(RDNA4)
262265
return ggml_cuda_mmq_get_config_rdna4(type, J, fallback);
266+
#elif defined(RDNA4_M)
267+
return ggml_cuda_mmq_get_config_rdna3_5(type, J, fallback);
263268
#elif defined(RDNA3_5)
264269
return ggml_cuda_mmq_get_config_rdna3_5(type, J, fallback);
265270
#elif defined(RDNA3)

‎ggml/src/ggml-cuda/mmvf.cu‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -838,7 +838,7 @@ bool ggml_cuda_should_use_mmvf(enum ggml_type type, int cc, const int64_t * src0
838838
return ne11 <= 8;
839839
} else if (GGML_CUDA_CC_IS_AMD(cc)) {
840840
if (fp16_mma_hardware_available(cc)) {
841-
if (GGML_CUDA_CC_IS_RDNA3(cc)) {
841+
if (GGML_CUDA_CC_IS_RDNA3(cc) || GGML_CUDA_CC_IS_RDNA4_M(cc)) {
842842
return ne11 <= 3;
843843
}
844844
if (GGML_CUDA_CC_IS_RDNA4(cc)) {

0 commit comments

Comments
 (0)