Repository navigation
Implement fp16 x fp8, bf16 x fp8, fp8 x fp16, and fp8 x bf16 for the V3 pipeline #4276
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.
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
Closed
assistant-librarian
wants to merge
15
commits into
develop
from
import/develop/ROCm_composable_kernel/pr-3669
Closed
Changes from all commits
Commits
Show all changes
15 commits
Select commit
Hold shift + click to select a range
90b950b
Rename the parameters of load_interleaved_pk_type and load_and_conver…
SamiAario-AMD 999b493
Add load_tile_transpose_convert for mixed precision transpose loading
SamiAario-AMD afba075
Add and use load_with_type_convert
SamiAario-AMD a9c5248
Add test_load_tile_transpose
SamiAario-AMD b381791
When defining c_tensor, replace hard coded values with values derived…
SamiAario-AMD 7a8c448
Add AttrNumAccessV_ as a template parameter to WarpGemmAttributeMfma:…
SamiAario-AMD 6ae8db7
Introduce DetermineWarpPrecType for determining warp GEMM precision t…
SamiAario-AMD 6fa3fde
Switch to an implementation of DetermineWarpPrecType that explicitly …
SamiAario-AMD 2e5fac9
Add MFMA warp gemm for float, float, float, 32, 32, 16
SamiAario-AMD 9475ac1
Add functionality and tests for bf16 x fp8 and fp8 x bf16
SamiAario-AMD abaa06b
Add functionality and tests for fp16 x fp8 and fp8 x fp16
SamiAario-AMD 3c4553a
Add type conversions to V4 pipeline, WIP!
SamiAario-AMD d168d2f
Refactor type conversions out of MakeBLdsBlockDescriptor, WIP!
SamiAario-AMD 1abad5a
Add a changelog entry
SamiAario-AMD 5a304ff
Add include statements added by remod.py
SamiAario-AMD File filter
Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
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
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
109 changes: 109 additions & 0 deletions
109
projects/composablekernel/include/ck_tile/ops/common/determine_warp_prec_type.hpp
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 |
|---|---|---|
| @@ -0,0 +1,109 @@ | ||
| // Copyright (c) Advanced Micro Devices, Inc., or its affiliates. | ||
| // SPDX-License-Identifier: MIT | ||
|
|
||
| #pragma once | ||
|
|
||
| #include "ck_tile/core.hpp" | ||
|
|
||
| // DetermineWarpPrecType is a set of rules to determine the right precision type to use | ||
| // for the warp GEMM, given the other precision type. This gives rise to a type conversion: | ||
| // type conversions are sometimes needed to obtain a pair of types that are compatible with | ||
| // the hardware matrix operations available. A typical use case is mixed precision GEMMs. | ||
|
|
||
| namespace ck_tile { | ||
| // For the most general case, default to no conversion. | ||
| template <typename APrecType, typename BPrecType> | ||
| struct DetermineWarpPrecType | ||
| { | ||
| using a_prec_type = APrecType; | ||
| using b_prec_type = BPrecType; | ||
| }; | ||
|
|
||
| // For pk_fp4_t x pk_fp4_t, keep pk_fp4_t | ||
| template <> | ||
| struct DetermineWarpPrecType<ck_tile::pk_fp4_t, ck_tile::pk_fp4_t> | ||
| { | ||
| using a_prec_type = ck_tile::pk_fp4_t; | ||
| using b_prec_type = ck_tile::pk_fp4_t; | ||
| }; | ||
|
|
||
| // For pk_int4_t x B, use the B type. | ||
| template <typename BPrecType> | ||
| struct DetermineWarpPrecType<ck_tile::pk_int4_t, BPrecType> | ||
| { | ||
| using a_prec_type = BPrecType; | ||
| using b_prec_type = BPrecType; | ||
| }; | ||
|
|
||
| // For A x pk_int4_t, use the A type. | ||
| template <typename APrecType> | ||
| struct DetermineWarpPrecType<APrecType, ck_tile::pk_int4_t> | ||
| { | ||
| using a_prec_type = APrecType; | ||
| using b_prec_type = APrecType; | ||
| }; | ||
|
|
||
| // For pk_fp4_t x B, use the B type. | ||
| template <typename BPrecType> | ||
| struct DetermineWarpPrecType<ck_tile::pk_fp4_t, BPrecType> | ||
| { | ||
| using a_prec_type = BPrecType; | ||
| using b_prec_type = BPrecType; | ||
| }; | ||
|
|
||
| // For A x pk_fp4_t, use the A type. | ||
| template <typename APrecType> | ||
| struct DetermineWarpPrecType<APrecType, ck_tile::pk_fp4_t> | ||
| { | ||
| using a_prec_type = APrecType; | ||
| using b_prec_type = APrecType; | ||
| }; | ||
|
|
||
| // For B x pk_fp4_raw_t, use the B type. | ||
| template <typename BPrecType> | ||
| struct DetermineWarpPrecType<ck_tile::pk_fp4_raw_t, BPrecType> | ||
| { | ||
| using a_prec_type = BPrecType; | ||
| using b_prec_type = BPrecType; | ||
| }; | ||
|
|
||
| // For A x pk_fp4_raw_t, use the A type. | ||
| template <typename APrecType> | ||
| struct DetermineWarpPrecType<APrecType, ck_tile::pk_fp4_raw_t> | ||
|
Comment on lines
+70
to
+72
|
||
| { | ||
| using a_prec_type = APrecType; | ||
| using b_prec_type = APrecType; | ||
| }; | ||
|
|
||
| // For fp8 x bf16, use fp8 | ||
| template <> | ||
| struct DetermineWarpPrecType<ck_tile::fp8_t, ck_tile::bf16_t> | ||
| { | ||
| using a_prec_type = ck_tile::fp8_t; | ||
| using b_prec_type = ck_tile::fp8_t; | ||
| }; | ||
|
|
||
| // For bf16 x fp8, use bf16 | ||
| template <> | ||
| struct DetermineWarpPrecType<ck_tile::bf16_t, ck_tile::fp8_t> | ||
| { | ||
| using a_prec_type = ck_tile::bf16_t; | ||
| using b_prec_type = ck_tile::bf16_t; | ||
| }; | ||
|
|
||
| // For fp8 x fp16, use fp8 | ||
| template <> | ||
| struct DetermineWarpPrecType<ck_tile::fp8_t, ck_tile::half_t> | ||
| { | ||
| using a_prec_type = ck_tile::fp8_t; | ||
| using b_prec_type = ck_tile::fp8_t; | ||
| }; | ||
|
|
||
| // For fp16 x fp8, use fp16 | ||
| template <> | ||
| struct DetermineWarpPrecType<ck_tile::half_t, ck_tile::fp8_t> | ||
| { | ||
| using a_prec_type = ck_tile::half_t; | ||
| using b_prec_type = ck_tile::half_t; | ||
| }; | ||
| }; // namespace ck_tile | ||
|
||
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
Oops, something went wrong.
Oops, something went wrong.
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.
The comments for the
pk_fp4_raw_tspecializations don’t match the template argument order (line 64 ispk_fp4_raw_t x B, notB x pk_fp4_raw_t). Updating these comments would reduce confusion when extending the rule set.