wmantly/llama.cpp-xrip
1
0 Fork 0
Code Issues Pull Requests Actions 32 Packages Projects Releases Wiki Activity

ROCm/HIP: add RDNA4_M family for gfx1170/gfx1171 (#31373)

Browse Source 
gfx117X is more like RDNA4, with 128b WMMA. But they're not identical. Adjust the defines accordingly to have a new RDNA4_M and use it where it makes sense.


Assisted-by: Codex
Assisted-by: GitHub Copilot

Co-authored-by: Mario Limonciello <mario.limonciello@amd.com>
This commit is contained in:
Slobodan JosicandMario Limonciello authored and GitHub committed 2026-10-29 18:23:29 +01:00
1 parent 68d20b884c
commit eb4e244be7
9 files changed
+72 -53
Download Patch File Download Diff File Expand all files Collapse all files

No files matched your search

ggml/src/ggml-cuda/common.cuh
+18 -8
View File
@@ -82,6 +82,7 @@
#define GGML_CUDA_CC_RDNA2 (GGML_CUDA_CC_OFFSET_AMD + 0x1030) // RX 6000, minimum for dp4a
#define GGML_CUDA_CC_RDNA3 (GGML_CUDA_CC_OFFSET_AMD + 0x1100) // RX 7000, minimum for WMMA
#define GGML_CUDA_CC_RDNA3_5 (GGML_CUDA_CC_OFFSET_AMD + 0x1150) // AI 370, AI Max 395 laptops.
#define GGML_CUDA_CC_RDNA4_M (GGML_CUDA_CC_OFFSET_AMD + 0x1170) // gfx117X
#define GGML_CUDA_CC_RDNA4 (GGML_CUDA_CC_OFFSET_AMD + 0x1200) // RX 9000
#define GGML_CUDA_CC_IS_AMD(cc) (cc >= GGML_CUDA_CC_OFFSET_AMD)
@@ -89,8 +90,9 @@
#define GGML_CUDA_CC_IS_RDNA1(cc) (cc >= GGML_CUDA_CC_RDNA1 && cc < GGML_CUDA_CC_RDNA2)
#define GGML_CUDA_CC_IS_RDNA2(cc) (cc >= GGML_CUDA_CC_RDNA2 && cc < GGML_CUDA_CC_RDNA3)
#define GGML_CUDA_CC_IS_RDNA3_0(cc) (cc >= GGML_CUDA_CC_RDNA3 && cc < GGML_CUDA_CC_RDNA3_5)
#define GGML_CUDA_CC_IS_RDNA3_5(cc) (cc >= GGML_CUDA_CC_RDNA3_5 && cc < GGML_CUDA_CC_RDNA4)
#define GGML_CUDA_CC_IS_RDNA3_5(cc) (cc >= GGML_CUDA_CC_RDNA3_5 && cc < GGML_CUDA_CC_RDNA4_M)
#define GGML_CUDA_CC_IS_RDNA3(cc) (GGML_CUDA_CC_IS_RDNA3_0(cc) || GGML_CUDA_CC_IS_RDNA3_5(cc))
#define GGML_CUDA_CC_IS_RDNA4_M(cc) (cc >= GGML_CUDA_CC_RDNA4_M && cc < GGML_CUDA_CC_RDNA4)
#define GGML_CUDA_CC_IS_RDNA4(cc) (cc >= GGML_CUDA_CC_RDNA4)
#define GGML_CUDA_CC_IS_GCN_APU(cc) ((cc) == GGML_CUDA_CC_GFX909 || (cc) == GGML_CUDA_CC_GFX90C)
#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))
@@ -277,9 +279,17 @@ static const char * cu_get_error_str(CUresult err) {
#define AMD_MFMA_AVAILABLE
#endif // defined(GGML_USE_HIP) && defined(CDNA) && !defined(GGML_HIP_NO_MMQ_MFMA)
#if defined(GGML_USE_HIP) && (defined(RDNA4) || defined(RDNA3))
#if defined(GGML_USE_HIP) && (defined(RDNA4) || defined(RDNA4_M))
#define AMD_WMMA_128B_AVAILABLE
#endif // defined(GGML_USE_HIP) && (defined(RDNA4) || defined(RDNA4_M))
#if defined(GGML_USE_HIP) && defined(RDNA3)
#define AMD_WMMA_256B_AVAILABLE
#endif // defined(GGML_USE_HIP) && defined(RDNA3)
#if defined(AMD_WMMA_128B_AVAILABLE) || defined(AMD_WMMA_256B_AVAILABLE)
#define AMD_WMMA_AVAILABLE
#endif // defined(GGML_USE_HIP) && defined(RDNA4)
#endif // defined(AMD_WMMA_128B_AVAILABLE) || defined(AMD_WMMA_256B_AVAILABLE)
// The Volta instructions are in principle available on Turing or newer but they are effectively unusable:
#if !defined(GGML_USE_HIP) && __CUDA_ARCH__ == GGML_CUDA_CC_VOLTA
@@ -326,7 +336,7 @@ static bool fast_fp16_hardware_available(const int cc) {
// To be used for feature selection of external libraries, e.g. cuBLAS.
static bool fp16_mma_hardware_available(const int cc) {
return (GGML_CUDA_CC_IS_NVIDIA(cc) && cc >= GGML_CUDA_CC_VOLTA) ||
GGML_CUDA_CC_IS_CDNA(cc) || GGML_CUDA_CC_IS_RDNA3(cc) || GGML_CUDA_CC_IS_RDNA4(cc) ||
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) ||
(GGML_CUDA_CC_IS_MTHREADS(cc) && cc >= GGML_CUDA_CC_QY2);
}
@@ -355,7 +365,7 @@ static bool amd_mfma_available(const int cc) {
}
static bool amd_wmma_available(const int cc) {
return (GGML_CUDA_CC_IS_RDNA4(cc) || GGML_CUDA_CC_IS_RDNA3(cc));
return GGML_CUDA_CC_IS_RDNA3(cc) || GGML_CUDA_CC_IS_RDNA4_M(cc) || GGML_CUDA_CC_IS_RDNA4(cc);
}
static bool volta_mma_available(const int cc) {
@@ -713,7 +723,7 @@ static __device__ __forceinline__ int ggml_cuda_dp4a(const int a, const int b, i
#if defined(GGML_USE_HIP)
#if defined(CDNA) || defined(RDNA2) || defined(__gfx906__)
c = __builtin_amdgcn_sdot4(a, b, c, false);
#elif defined(RDNA3) || defined(RDNA4)
#elif defined(RDNA3) || defined(RDNA4_M) || defined(RDNA4)
c = __builtin_amdgcn_sudot4( true, a, true, b, c, false);
#elif defined(RDNA1) || defined(__gfx900__)
int tmp1;
@@ -758,9 +768,9 @@ static __device__ __forceinline__ void ggml_cuda_mad(float & acc, const float2 v
acc += v.y*u.y;
}
#if defined(GGML_USE_HIP) && (defined(RDNA2) || defined(RDNA3) || defined(RDNA4) || defined(__gfx906__) || defined(CDNA))
#if defined(GGML_USE_HIP) && (defined(RDNA2) || defined(RDNA3) || defined(RDNA4_M) || defined(RDNA4) || defined(__gfx906__) || defined(CDNA))
#define V_DOT2_F32_F16_AVAILABLE
#endif // defined(GGML_USE_HIP) && (defined(RDNA2) || defined(RDNA3) || defined(RDNA4) || defined(__gfx906__) || defined(CDNA))
#endif // defined(GGML_USE_HIP) && (defined(RDNA2) || defined(RDNA3) || defined(RDNA4_M) || defined(RDNA4) || defined(__gfx906__) || defined(CDNA))
static __device__ __forceinline__ void ggml_cuda_mad(float & acc, const half2 v, const half2 u) {
#ifdef V_DOT2_F32_F16_AVAILABLE
ggml/src/ggml-cuda/fattn-mma-f16.cuh
+1 -1
View File
@@ -1255,7 +1255,7 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile(
T_B_KQ Q_B[(Q_in_reg ? DKQ/(2*T_B_KQ::J) : 1)];
#if defined(TURING_MMA_AVAILABLE)
T_C_VKQ VKQ_C[cols_per_warp == 8 ? DV/T_C_VKQ::I : DV/(2*T_C_VKQ::J)];
#elif defined(AMD_WMMA_AVAILABLE) && defined(RDNA3)
#elif defined(AMD_WMMA_256B_AVAILABLE)
T_C_VKQ VKQ_C[DV % 32 != 0 ? DV/T_C_VKQ::J : DV/(2*T_C_VKQ::J)];
#elif defined(AMD_MFMA_AVAILABLE)
T_C_VKQ VKQ_C[ DV/T_C_VKQ::J];
ggml/src/ggml-cuda/ggml-cuda.cu
+1 -1
View File
@@ -1565,7 +1565,7 @@ static void ggml_cuda_mul_mat_cublas_impl(ggml_backend_cuda_context & ctx, const
if (compute_type == GGML_TYPE_F16) {
prefer_f32_output = cc == GGML_CUDA_CC_VOLTA || GGML_CUDA_CC_IS_RDNA4(cc) || GGML_CUDA_CC_IS_CDNA(cc);
} else if (compute_type == GGML_TYPE_BF16) {
prefer_f32_output = !GGML_CUDA_CC_IS_RDNA3(cc) && !GGML_CUDA_CC_IS_CDNA(cc);
prefer_f32_output = !GGML_CUDA_CC_IS_RDNA3(cc) && !GGML_CUDA_CC_IS_RDNA4_M(cc) && !GGML_CUDA_CC_IS_CDNA(cc);
}
if (prefer_f32_output) {
ggml/src/ggml-cuda/mma.cuh
+31 -31
View File
@@ -76,11 +76,11 @@ namespace ggml_cuda_mma {
// For the A/C matrices this means I major == row major, J major == column major.
// For the B matrix this means I major == column major, J major == row major.
// MIRRORED == Each data value is held exactly once per thread subgroup.
DATA_LAYOUT_I_MAJOR = 0, // Always used for Turing, Ampere, Ada Lovelace, consumer Blackwell, matrix A&B for RDNA4 and CDNA.
DATA_LAYOUT_J_MAJOR = 10, // Matrix C for CDNA and RDNA4, int and float matrix C for RDNA3.
DATA_LAYOUT_I_MAJOR = 0, // Always used for Turing, Ampere, Ada Lovelace, consumer Blackwell, matrix A&B for WMMA128 and CDNA.
DATA_LAYOUT_J_MAJOR = 10, // Matrix C for CDNA and WMMA128, int and float matrix C for WMMA256.
DATA_LAYOUT_I_MAJOR_MIRRORED = 20, // Volta, matrix A&B for RDNA3.
DATA_LAYOUT_J_MAJOR_MIRRORED = 30,
DATA_LAYOUT_I_MAJOR_SCRAMBLED = 40, // Scrambled matrix C for faster transposition (RDNA4/CDNA), convert to float to unscramble.
DATA_LAYOUT_I_MAJOR_SCRAMBLED = 40, // Scrambled matrix C for faster transposition (WMMA128/CDNA), convert to float to unscramble.
};
// Implemented mma combinations are:
// - (I_MAJOR, I_MAJOR) -> I_MAJOR
@@ -88,11 +88,11 @@ namespace ggml_cuda_mma {
// - (I_MAJOR, J_MAJOR_MIRRORED) -> I_MAJOR
static constexpr __device__ data_layout get_input_data_layout() {
#if defined(RDNA3) || defined(VOLTA_MMA_AVAILABLE)
#if defined(AMD_WMMA_256B_AVAILABLE) || defined(VOLTA_MMA_AVAILABLE)
return DATA_LAYOUT_I_MAJOR_MIRRORED;
#else
return DATA_LAYOUT_I_MAJOR;
#endif // defined(RDNA3) || defined(VOLTA_MMA_AVAILABLE)
#endif // defined(AMD_WMMA_256B_AVAILABLE) || defined(VOLTA_MMA_AVAILABLE)
}
template <int I_, int J_, typename T, data_layout ds_=DATA_LAYOUT_I_MAJOR>
@@ -201,7 +201,7 @@ namespace ggml_cuda_mma {
static __device__ __forceinline__ int get_j(const int l) {
if constexpr (I == 16 && J == 16) {
#if defined(RDNA3)
#if defined(AMD_WMMA_256B_AVAILABLE)
if constexpr (std::is_same_v<T, float> || std::is_same_v<T, int>) {
// matrix C
return 2 * l + (threadIdx.x / 16);
@@ -210,11 +210,11 @@ namespace ggml_cuda_mma {
return l;
}
#else
// matrix C is the transposed matrix A&B on RDNA4
// matrix C is the transposed matrix A&B on WMMA128
return ne * (threadIdx.x / 16) + l;
#endif // defined(RDNA3)
#endif // defined(AMD_WMMA_256B_AVAILABLE)
} else if constexpr (I == 16 && J == 8) {
// mmq input for RDNA4
// mmq input for WMMA128
return ne * (threadIdx.x / 16) + l;
} else if constexpr (I == 16 && J == 4) {
return ne * (threadIdx.x / 16) + l;
@@ -573,7 +573,7 @@ namespace ggml_cuda_mma {
static constexpr int I = I_;
static constexpr int J = J_;
static constexpr data_layout dl = DATA_LAYOUT_I_MAJOR_MIRRORED;
#if defined(RDNA3)
#if defined(AMD_WMMA_256B_AVAILABLE)
static constexpr int ne = tile<I_, J_, float, DATA_LAYOUT_I_MAJOR_MIRRORED>::ne;
half2 x[ne] = {{0.0f, 0.0f}};
@@ -616,7 +616,7 @@ namespace ggml_cuda_mma {
return -1;
}
}
#endif // defined(RDNA3)
#endif // defined(AMD_WMMA_256B_AVAILABLE)
};
template <int I_, int J_>
@@ -694,7 +694,7 @@ namespace ggml_cuda_mma {
};
static __device__ __forceinline__ tile<16, 16, half2, DATA_LAYOUT_I_MAJOR> unscramble(const tile<16, 16, half2, DATA_LAYOUT_I_MAJOR_SCRAMBLED> & t) {
#if defined(AMD_MFMA_AVAILABLE) || (defined(AMD_WMMA_AVAILABLE) && defined(RDNA4))
#if defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_128B_AVAILABLE)
tile<16, 16, half2, DATA_LAYOUT_I_MAJOR> ret;
#pragma unroll
for (int l0 = 0; l0 < t.ne/2; ++l0) {
@@ -705,7 +705,7 @@ namespace ggml_cuda_mma {
#else
NO_DEVICE_CODE;
GGML_UNUSED(t);
#endif // defined(AMD_MFMA_AVAILABLE) || (defined(AMD_WMMA_AVAILABLE) && defined(RDNA4))
#endif // defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_128B_AVAILABLE)
}
#if defined(TURING_MMA_AVAILABLE)
@@ -726,7 +726,7 @@ namespace ggml_cuda_mma {
return ret;
}
#elif defined(AMD_WMMA_AVAILABLE) && defined(RDNA3)
#elif defined(AMD_WMMA_256B_AVAILABLE)
static __device__ __forceinline__ tile<16, 8, half2, DATA_LAYOUT_I_MAJOR_MIRRORED> get_half2(
const tile<16, 16, float, DATA_LAYOUT_I_MAJOR> & tile_float) {
tile<16, 8, half2, DATA_LAYOUT_I_MAJOR_MIRRORED> ret;
@@ -1171,7 +1171,7 @@ namespace ggml_cuda_mma {
: "r"(Axi[2]), "r"(Axi[3]), "r"(Bxi[3]));
#endif // __CUDA_ARCH__ >= GGML_CUDA_CC_AMPERE
#elif defined(AMD_WMMA_AVAILABLE)
#if defined(RDNA4)
#if defined(AMD_WMMA_128B_AVAILABLE)
using halfx8_t = __attribute__((ext_vector_type(8))) _Float16;
halfx8_t& acc_frag = reinterpret_cast<halfx8_t&>(D.x[0]);
const halfx8_t& a_frag = reinterpret_cast<const halfx8_t&>(A.x[0]);
@@ -1180,7 +1180,7 @@ namespace ggml_cuda_mma {
#else
GGML_UNUSED_VARS(D, A, B);
NO_DEVICE_CODE;
#endif // defined(RDNA4)
#endif // defined(AMD_WMMA_128B_AVAILABLE)
#elif defined(AMD_MFMA_AVAILABLE)
// MFMA: FP16 input, FP32 accumulate, convert back to half2.
using halfx4_t = __attribute__((ext_vector_type(4))) _Float16;
@@ -1218,7 +1218,7 @@ namespace ggml_cuda_mma {
static __device__ __forceinline__ void mma(
tile<16, 16, half2, DATA_LAYOUT_I_MAJOR_SCRAMBLED> & D, const tile<32, 8, half2, DATA_LAYOUT_I_MAJOR> & A,
const tile<16, 8, half2, DATA_LAYOUT_I_MAJOR> & B) {
#if defined(AMD_MFMA_AVAILABLE) || (defined(AMD_WMMA_AVAILABLE) && defined(RDNA4))
#if defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_128B_AVAILABLE)
tile<16, 8, half2> * D16 = (tile<16, 8, half2> *) &D;
const tile<16, 8, half2> * A16 = (const tile<16, 8, half2> *) &A;
mma(D16[0], A16[0], B);
@@ -1226,7 +1226,7 @@ namespace ggml_cuda_mma {
#else
GGML_UNUSED_VARS(D, A, B);
NO_DEVICE_CODE;
#endif // defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) && defined(RDNA4)
#endif // defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_128B_AVAILABLE)
}
template <data_layout dl_ab, data_layout dl_d>
@@ -1373,14 +1373,14 @@ namespace ggml_cuda_mma {
: "r"(Axi[2]), "r"(Axi[3]), "r"(Bxi[3]));
#endif // __CUDA_ARCH__ >= GGML_CUDA_CC_AMPERE
#elif defined(AMD_WMMA_AVAILABLE)
#if defined(RDNA4)
#if defined(AMD_WMMA_128B_AVAILABLE)
using halfx8_t = __attribute__((ext_vector_type(8))) _Float16;
using floatx8_t = __attribute__((ext_vector_type(8))) float;
floatx8_t& acc_frag = reinterpret_cast<floatx8_t&>(D.x[0]);
const halfx8_t& a_frag = reinterpret_cast<const halfx8_t&>(A.x[0]);
const halfx8_t& b_frag = reinterpret_cast<const halfx8_t&>(B.x[0]);
acc_frag = __builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12(a_frag, b_frag, acc_frag);
#elif defined(RDNA3)
#elif defined(AMD_WMMA_256B_AVAILABLE)
using halfx16_t = __attribute__((ext_vector_type(16))) _Float16;
using floatx8_t = __attribute__((ext_vector_type(8))) float;
floatx8_t& acc_frag = reinterpret_cast<floatx8_t&>(D.x[0]);
@@ -1390,7 +1390,7 @@ namespace ggml_cuda_mma {
#else
GGML_UNUSED_VARS(D, A, B);
NO_DEVICE_CODE;
#endif // RDNA4
#endif // AMD_WMMA_128B_AVAILABLE
#elif defined(AMD_MFMA_AVAILABLE)
using halfx4_t = __attribute__((ext_vector_type(4))) _Float16;
using floatx4_t = __attribute__((ext_vector_type(4))) float;
@@ -1408,14 +1408,14 @@ namespace ggml_cuda_mma {
static __device__ __forceinline__ void mma(
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) {
#if defined(AMD_WMMA_AVAILABLE)
#if defined(RDNA4)
#if defined(AMD_WMMA_128B_AVAILABLE)
using bf16x8_t = __attribute__((ext_vector_type(8))) short;
using floatx8_t = __attribute__((ext_vector_type(8))) float;
floatx8_t& acc_frag = reinterpret_cast<floatx8_t&>(D.x[0]);
const bf16x8_t& a_frag = reinterpret_cast<const bf16x8_t&>(A.x[0]);
const bf16x8_t& b_frag = reinterpret_cast<const bf16x8_t&>(B.x[0]);
acc_frag = __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32_gfx12(a_frag, b_frag, acc_frag);
#elif defined(RDNA3)
#elif defined(AMD_WMMA_256B_AVAILABLE)
using bf16x16_t = __attribute__((ext_vector_type(16))) short;
using floatx8_t = __attribute__((ext_vector_type(8))) float;
floatx8_t& acc_frag = reinterpret_cast<floatx8_t&>(D.x[0]);
@@ -1425,7 +1425,7 @@ namespace ggml_cuda_mma {
#else
GGML_UNUSED_VARS(D, A, B);
NO_DEVICE_CODE;
#endif // defined(RDNA4)
#endif // defined(AMD_WMMA_128B_AVAILABLE)
#elif defined(AMD_MFMA_AVAILABLE)
using floatx4_t = __attribute__((ext_vector_type(4))) float;
floatx4_t& acc_frag = reinterpret_cast<floatx4_t&>(D.x[0]);
@@ -1467,19 +1467,19 @@ namespace ggml_cuda_mma {
#elif defined(AMD_WMMA_AVAILABLE)
using int32x8_t = __attribute__((__vector_size__(8 * sizeof(int)))) int;
int32x8_t * acc = (int32x8_t *) D.x;
#if defined(RDNA4)
#if defined(AMD_WMMA_128B_AVAILABLE)
using int32x2_t = __attribute__((__vector_size__(2 * sizeof(int)))) int;
int32x2_t * a_vec = (int32x2_t *) A.x;
int32x2_t * b_vec = (int32x2_t *) B.x;
acc[0] = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12(true, a_vec[0], true, b_vec[0], acc[0], true);
acc[0] = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12(true, a_vec[1], true, b_vec[1], acc[0], true);
#elif defined(RDNA3)
#elif defined(AMD_WMMA_256B_AVAILABLE)
using int32x4_t = __attribute__((__vector_size__(4 * sizeof(int)))) int;
int32x4_t * a_vec = (int32x4_t *) A.x;
int32x4_t * b_vec = (int32x4_t *) B.x;
acc[0] = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32(true, a_vec[0], true, b_vec[0], acc[0], true);
acc[0] = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32(true, a_vec[1], true, b_vec[1], acc[0], true);
#endif // RDNA4
#endif // AMD_WMMA_128B_AVAILABLE
#else
GGML_UNUSED_VARS(D, A, B);
NO_DEVICE_CODE;
@@ -1556,7 +1556,7 @@ namespace ggml_cuda_mma {
static __device__ __forceinline__ void mma(
tile<16, 16, half2, DATA_LAYOUT_I_MAJOR> & D, const tile<32, 8, half2, DATA_LAYOUT_I_MAJOR_MIRRORED> & A,
const tile<16, 8, half2, DATA_LAYOUT_I_MAJOR_MIRRORED> & B) {
#if defined(AMD_WMMA_AVAILABLE) && defined(RDNA3)
#if defined(AMD_WMMA_256B_AVAILABLE)
using halfx16_t = __attribute__((ext_vector_type(16))) _Float16;
halfx16_t * xD = (halfx16_t *) D.x;
const halfx16_t * xA = (const halfx16_t *) A.x;
@@ -1585,17 +1585,17 @@ namespace ggml_cuda_mma {
#elif defined(AMD_WMMA_AVAILABLE)
using int32x8_t = __attribute__((__vector_size__(8 * sizeof(int)))) int;
int32x8_t * acc = (int32x8_t *) D.x;
#if defined(RDNA4)
#if defined(AMD_WMMA_128B_AVAILABLE)
using int32x2_t = __attribute__((__vector_size__(2 * sizeof(int)))) int;
int32x2_t * a_vec = (int32x2_t *) A.x;
int32x2_t * b_vec = (int32x2_t *) B.x;
acc[0] = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12(true, a_vec[0], true, b_vec[0], acc[0], false);
#elif defined(RDNA3)
#elif defined(AMD_WMMA_256B_AVAILABLE)
using int32x4_t = __attribute__((__vector_size__(4 * sizeof(int)))) int;
int32x4_t * a_vec = (int32x4_t *) A.x;
int32x4_t * b_vec = (int32x4_t *) B.x;
acc[0] = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32(true, a_vec[0], true, b_vec[0], acc[0], false);
#endif // RDNA4
#endif // AMD_WMMA_128B_AVAILABLE
#else
GGML_UNUSED(D);
GGML_UNUSED(A);
ggml/src/ggml-cuda/mmq.cu
+3 -3
View File
@@ -242,7 +242,7 @@ void ggml_cuda_mul_mat_q(
// Each expert only sees ne12*n_expert_used/ne02 tokens on average.
// On RDNA3 and RDNA4 it is faster to pick the tile size against this value instead of ne12.
if (GGML_CUDA_CC_IS_RDNA3(cc) || GGML_CUDA_CC_IS_RDNA4(cc)) {
if (GGML_CUDA_CC_IS_RDNA3(cc) || GGML_CUDA_CC_IS_RDNA4_M(cc) || GGML_CUDA_CC_IS_RDNA4(cc)) {
ncols_opt = (ne12*n_expert_used + ne02 - 1) / ne02;
}
}
@@ -470,7 +470,7 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t
}
if (amd_wmma_available(cc)) {
if (GGML_CUDA_CC_IS_RDNA3(cc)) {
if (GGML_CUDA_CC_IS_RDNA3(cc) || GGML_CUDA_CC_IS_RDNA4_M(cc)) {
// High expert counts are almost always better on MMQ due to
// the synchronization overhead in the cuBLAS/hipBLAS path:
// https://github.com/ggml-org/llama.cpp/pull/18202
@@ -487,7 +487,7 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t
return ne11 <= (GGML_CUDA_CC_IS_RDNA3_0(cc) ? 128 : 256);
case GGML_TYPE_IQ2_XS:
case GGML_TYPE_IQ2_S:
return GGML_CUDA_CC_IS_RDNA3_5(cc) || ne11 <= 128;
return GGML_CUDA_CC_IS_RDNA3_5(cc) || GGML_CUDA_CC_IS_RDNA4_M(cc) || ne11 <= 128;
default:
return true;
}
ggml/src/ggml-cuda/mmq.cuh
+5
View File
@@ -240,6 +240,9 @@ static __host__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config(const ggml_type ty
if (GGML_CUDA_CC_IS_RDNA4(cc)) {
return ggml_cuda_mmq_get_config_rdna4(type, J, fallback);
}
if (GGML_CUDA_CC_IS_RDNA4_M(cc)) {
return ggml_cuda_mmq_get_config_rdna3_5(type, J, fallback);
}
if (GGML_CUDA_CC_IS_RDNA3_5(cc)) {
return ggml_cuda_mmq_get_config_rdna3_5(type, J, fallback);
}
@@ -272,6 +275,8 @@ static constexpr __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config(ggml_t
return ggml_cuda_mmq_get_config_cdna(type, J, fallback);
#elif defined(RDNA4)
return ggml_cuda_mmq_get_config_rdna4(type, J, fallback);
#elif defined(RDNA4_M)
return ggml_cuda_mmq_get_config_rdna3_5(type, J, fallback);
#elif defined(RDNA3_5)
return ggml_cuda_mmq_get_config_rdna3_5(type, J, fallback);
#elif defined(RDNA3)
ggml/src/ggml-cuda/mmvf.cu
+1 -1
View File
@@ -844,7 +844,7 @@ bool ggml_cuda_should_use_mmvf(enum ggml_type type, int cc, int warp_size, const
return ne11 <= MMVF_MAX_BATCH_SIZE;
}
if (fp16_mma_hardware_available(cc)) {
if (GGML_CUDA_CC_IS_RDNA3(cc)) {
if (GGML_CUDA_CC_IS_RDNA3(cc) || GGML_CUDA_CC_IS_RDNA4_M(cc)) {
return ne11 <= 3;
}
if (GGML_CUDA_CC_IS_RDNA4(cc)) {
ggml/src/ggml-cuda/mmvq.cu
+4 -4
View File
@@ -108,7 +108,7 @@ static constexpr __device__ mmvq_parameter_table_id get_device_table_id() {
return MMVQ_PARAMETERS_RDNA4;
#elif defined(RDNA3_0)
return MMVQ_PARAMETERS_RDNA3_0;
#elif defined(RDNA2) || defined(RDNA3_5)
#elif defined(RDNA2) || defined(RDNA3_5) || defined(RDNA4_M)
return MMVQ_PARAMETERS_RDNA2;
#elif defined(GCN) || defined(CDNA)
return MMVQ_PARAMETERS_GCN;
@@ -128,7 +128,7 @@ static __host__ mmvq_parameter_table_id get_device_table_id(int cc) {
if (GGML_CUDA_CC_IS_RDNA3_0(cc)) {
return MMVQ_PARAMETERS_RDNA3_0;
}
if (GGML_CUDA_CC_IS_RDNA2(cc) || GGML_CUDA_CC_IS_RDNA3_5(cc)) {
if (GGML_CUDA_CC_IS_RDNA2(cc) || GGML_CUDA_CC_IS_RDNA3_5(cc) || GGML_CUDA_CC_IS_RDNA4_M(cc)) {
return MMVQ_PARAMETERS_RDNA2;
}
if (GGML_CUDA_CC_IS_GCN(cc) || GGML_CUDA_CC_IS_CDNA(cc)) {
@@ -299,7 +299,7 @@ int get_mmvq_mmid_max_batch(ggml_type type, int cc) {
if (GGML_CUDA_CC_IS_RDNA4(cc)) {
return get_mmvq_mmid_max_batch_rdna4(type);
}
if (GGML_CUDA_CC_IS_RDNA3(cc)) {
if (GGML_CUDA_CC_IS_RDNA3(cc) || GGML_CUDA_CC_IS_RDNA4_M(cc)) {
return get_mmvq_mmid_max_batch_rdna3(type);
}
if (GGML_CUDA_CC_IS_RDNA1(cc) || GGML_CUDA_CC_IS_RDNA2(cc)) {
@@ -432,7 +432,7 @@ template <ggml_type type>
static constexpr __device__ int get_mmvq_mmid_max_batch_for_device() {
#if defined(RDNA4)
return get_mmvq_mmid_max_batch_rdna4(type);
#elif defined(RDNA3)
#elif defined(RDNA3) || defined(RDNA4_M)
return get_mmvq_mmid_max_batch_rdna3(type);
#elif defined(RDNA2) || defined(RDNA1)
return get_mmvq_mmid_max_batch_rdna1_rdna2(type);
ggml/src/ggml-cuda/vendors/hip.h
Vendored
+8 -4
View File
@@ -216,9 +216,13 @@
#define RDNA4
#endif // defined(__GFX12__)
#if defined(__GFX11__)
#if defined(__gfx1170__) || defined(__gfx1171__)
#define RDNA4_M
#endif // defined(__gfx1170__) || defined(__gfx1171__)
#if defined(__GFX11__) && !defined(RDNA4_M)
#define RDNA3
#endif // defined(__GFX11__)
#endif // defined(__GFX11__) && !defined(RDNA4_M)
#if defined(__gfx1150__) || defined(__gfx1151__) || defined(__gfx1152__) || defined(__gfx1153__)
#define RDNA3_5
@@ -237,9 +241,9 @@
#define RDNA1
#endif // defined(__gfx1010__) || defined(__gfx1012__)
#if defined(RDNA4) || defined(RDNA3) || defined(RDNA2) || defined(RDNA1)
#if defined(RDNA4) || defined(RDNA4_M) || defined(RDNA3) || defined(RDNA2) || defined(RDNA1)
#define RDNA // For the entire family
#endif // defined(RDNA4) || defined(RDNA3) || defined(RDNA2) || defined(RDNA1)
#endif // defined(RDNA4) || defined(RDNA4_M) || defined(RDNA3) || defined(RDNA2) || defined(RDNA1)
#ifndef __has_builtin
#define __has_builtin(x) 0