HIP: Add RDNA3 WMMA support to MMF #17495
+130
−43
There are no files selected for viewing
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
|
|
@@ -151,7 +151,7 @@ bool ggml_cuda_should_use_mmf(enum ggml_type type, int cc, int warp_size, const | |
| return false; | ||
| } | ||
| } else { | ||
| if (src1_ncols > 16 || amd_wmma_available(cc)) { | ||
| return false; | ||
| } | ||
| } | ||
|
|
||
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
|
|
@@ -306,11 +306,11 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11) { | |
| return false; | ||
| } | ||
|
|
||
| // RDNA4 has integer WMMA support, always use MMQ | ||
| if (GGML_CUDA_CC_IS_RDNA4(cc)) { | ||
| return true; | ||
| } | ||
|
|
||
| // RDNA3 uses DP4A path (no integer WMMA), enable for small batches | ||
| return (!GGML_CUDA_CC_IS_RDNA3(cc) && !GGML_CUDA_CC_IS_CDNA(cc)) || ne11 < MMQ_DP4A_MAX_BATCH_SIZE; | ||
| } | ||
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
|
|
@@ -256,7 +256,7 @@ static constexpr __device__ int mmq_get_granularity_device(const int /*mmq_x*/) | |
|
|
||
| #if defined(GGML_USE_HIP) | ||
| static int mmq_get_nwarps_host(const int cc, const int warp_size) { | ||
| return (amd_mfma_available(cc) || amd_wmma_available(cc)) ? 8 : 256/warp_size; | ||
|
Collaborator
There was a problem hiding this comment. Why are you changing this? |
||
| } | ||
| #else | ||
| static int mmq_get_nwarps_host(const int /*cc*/, const int warp_size) { | ||
|
|
@@ -1178,7 +1178,7 @@ static __device__ __forceinline__ void vec_dot_q8_0_16_q8_1_mma( | |
| } | ||
| } | ||
| } | ||
| #elif defined(AMD_WMMA_INT_AVAILABLE) //wmma instructions can handle 16x4 tiles, does not require loading 64x2 tiles | ||
| typedef tile<16, 4, int> tile_A; | ||
| typedef tile<16, 4, int> tile_B; | ||
| typedef tile<16, 16, int> tile_C; | ||
|
|
@@ -1500,7 +1500,7 @@ static __device__ __forceinline__ void vec_dot_q2_K_q8_1_mma( | |
| } | ||
| } | ||
| } | ||
| #elif defined(AMD_WMMA_INT_AVAILABLE) //wmma instructions can handle 16x4 tiles, does not require loading 64x2 tiles | ||
|
|
||
| typedef tile<16, 4, int> tile_A; | ||
| typedef tile<16, 4, int> tile_B; | ||
|
|
@@ -1888,7 +1888,7 @@ template <int mmq_y, bool need_check> static __device__ __forceinline__ void loa | |
| #else | ||
| int i = (i0 + threadIdx.y*rows_per_warp + threadIdx.x/2) % mmq_y; | ||
| { | ||
| #endif // defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_INT_AVAILABLE) | ||
| if (need_check) { | ||
| i = min(i, i_max); | ||
| } | ||
|
|
@@ -2313,7 +2313,7 @@ static __device__ __forceinline__ void vec_dot_q6_K_q8_1_mma( | |
| } | ||
| } | ||
| } | ||
| #elif defined(AMD_WMMA_INT_AVAILABLE) //wmma instructions can handle 16x4 tiles, does not require loading 64x2 tiles | ||
| typedef tile<16, 4, int> tile_A; | ||
| typedef tile<16, 4, int> tile_B; | ||
| typedef tile<16, 16, int> tile_C; | ||
|
|
@@ -3018,7 +3018,7 @@ static __device__ __forceinline__ void mmq_write_back_mma( | |
| #else | ||
| typedef tile<16, 8, int> tile_C; | ||
| constexpr int rows_per_warp = 2 * granularity; | ||
| #endif // defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_INT_AVAILABLE) | ||
| constexpr int ntx = rows_per_warp/tile_C::I; // Number of x minitiles per warp. | ||
|
|
||
| const int i0 = (threadIdx.y / ntx) * (ntx*tile_C::I); | ||
|
|
@@ -3232,7 +3232,7 @@ static __device__ __forceinline__ void mul_mat_q_process_tile( | |
| #else | ||
| constexpr vec_dot_mmq_t vec_dot = mmq_type_traits<mmq_x, mmq_y, need_check, type>::vec_dot_dp4a; | ||
| constexpr mmq_write_back_t write_back = mmq_write_back_dp4a<mmq_x, mmq_y, need_check>; | ||
| #endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_INT_AVAILABLE) | ||
|
|
||
| constexpr int blocks_per_iter = MMQ_ITER_K / qk; | ||
|
|
||
|
|
||
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
As I said before, use the term "not implemented" instead of "not supported" when talking about RDNA3 int8 WMMA in the context of llama.cpp/ggml.