Cleanup and refactoring related to tile loading #3505
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
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 |
|---|---|---|
|
|
@@ -3,48 +3,43 @@ | |
|
|
||
| #pragma once | ||
|
|
||
| #include "ck_tile/core.hpp" | ||
| #include "ck_tile/ops/elementwise/unary_element_wise_operation.hpp" | ||
|
|
||
| namespace ck_tile { | ||
|
|
||
| template <typename DstDataType, index_t UnaryOpSize> | ||
| struct ConverterLoader | ||
| { | ||
| template <typename WarpWindow, typename WarpTile> | ||
| CK_TILE_DEVICE static void load_interleaved_pk_type(WarpTile& dst, const WarpWindow& src) | ||
| { | ||
| static_assert(WarpTile::get_thread_buffer_size() % UnaryOpSize == 0); | ||
| constexpr index_t thread_buffer_size = WarpTile::get_thread_buffer_size() / UnaryOpSize; | ||
| const auto tmp = load_tile(src); | ||
|
|
||
| using DstVectorType = DstDataType __attribute__((ext_vector_type(UnaryOpSize))); | ||
| static_for<0, thread_buffer_size, 1>{}([&](auto i) { | ||
| const element_wise::PassThroughPack8 elementwise_op{}; | ||
|
|
||
| elementwise_op(dst.get_thread_buffer().template get_as<DstVectorType>()(i), | ||
| tmp.get_thread_buffer().template get_as<pk_int4x4_t>()[i]); | ||
| }); | ||
| } | ||
| }; | ||
|
|
||
| template <index_t UnaryOpSize, bool LoadTranspose = false, typename WarpTile, typename WarpWindow> | ||
| CK_TILE_DEVICE void load_and_convert_tile(WarpTile& dst, const WarpWindow& src) | ||
joyeamd marked this conversation as resolved.
Show resolved
Hide resolved
|
||
| { | ||
| if constexpr(std::is_same_v<typename WarpWindow::Base::DataType, pk_int4_t>) | ||
| { | ||
| static_assert(!LoadTranspose, "LoadTranspose not supported with pk_int4_t"); | ||
| ConverterLoader<typename WarpTile::DataType, UnaryOpSize>::load_interleaved_pk_type(dst, | ||
| src); | ||
| } | ||
| else if constexpr(LoadTranspose) | ||
| { | ||
| load_tile_transpose(dst, src); | ||
|
Contributor
Author
There was a problem hiding this comment. Copilot seems to agree with me that assignment cannot be optimized by the compiler in such a way that the creation of temporaries is avoided, when assigning to a complex type. It then ought to make sense to avoid assignment when possible, and pass the object to be assigned to as a reference instead.
Collaborator
There was a problem hiding this comment. @SamiAario-AMD Have you checked the assembly for this change? Does the register usage is different? |
||
| } | ||
| else | ||
| { | ||
|
|
||
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
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.
tmp is not the best name. What about naming it
srcand the windowsrc_window?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.
I will do this in a separate PR as discussed.