[CK_BUILDER] Add grouped conv fwd ck tile profiler #3518
+3,083
−65
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
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,91 @@ | ||
| // Copyright (c) Advanced Micro Devices, Inc., or its affiliates. | ||
| // SPDX-License-Identifier: MIT | ||
|
|
||
| #pragma once | ||
|
|
||
| #include "ck_tile/builder/testing/conv_fwd.hpp" | ||
| #include "ck_tile/host/kernel_launch.hpp" | ||
| #include "ck_tile/builder/factory/helpers/ck/conv_elementwise_op.hpp" | ||
| #include "ck_tile/ops/gemm.hpp" | ||
| #include "ck_tile/ops/grouped_convolution.hpp" | ||
| #include <type_traits> | ||
| #include <array> | ||
|
|
||
| /// This file contains the implementation details for invoking/testing | ||
| /// grouped convolution operations in CK Tile. The main item is the | ||
| /// `run()` function, which is the main implementation used to invoke | ||
| /// CK Tile grouped forward convolution kernels. | ||
|
|
||
| namespace ck_tile::builder::test { | ||
|
|
||
| namespace detail { | ||
|
|
||
| /// @brief Concept for checking whether this is the CK Tile convolution | ||
| /// implementation. | ||
| /// | ||
| /// This is the same as `::ck_tile::builder::test::CkConvInstance`, except | ||
| /// with some utility aliases. For that reason, its moved to this detail | ||
| /// namespace. | ||
| template <typename Conv, auto SIGNATURE> | ||
| concept CkTileConvInstance = requires(Conv&) { | ||
| { Conv::BlockSize() }; | ||
| }; | ||
|
|
||
| } // namespace detail | ||
|
|
||
| /// @brief Concept for checking whether a convolution is invoked like CK Tile. | ||
|
Collaborator
There was a problem hiding this comment. Note: we no longer need the doxygen comments. AMD is not running this code through doxygen for the CK / CK Tile API. Keep the file comments, but you don't need the triple slashes or the doxygen tags.
Contributor
Author
There was a problem hiding this comment. Even if we dont use doxygen I think it looks nice and it is commonly used in other headers. Let me know if you prefer to change it if no I would like to keep this
Collaborator
There was a problem hiding this comment. Eventually we'll probably strip out the doxygen tags, but I'm fine with it either way now. I'm just letting you know you don't have to use them. |
||
| /// | ||
| /// This concept is used to tell whether a convolution implementation is | ||
| /// likely to be an "CK Tile" implementation - that is, whether we should | ||
| /// invoke it as an CK Tile kernel. This is mainly used with `run()` to | ||
| /// differentiate which implementation that should be invoked. | ||
| /// | ||
| /// - SIGNATURE is the operation signature. | ||
| /// - Conv is a convolution instance created by the CK Builder API. | ||
| template <typename Conv, auto SIGNATURE> | ||
| concept CkTileConvInstance = detail::CkTileConvInstance<Conv, SIGNATURE>; | ||
|
|
||
| /// @brief `run()` specialization for forward convolution and CK Tile. | ||
| /// | ||
| /// @tparam SIGNATURE Forward convolution signature. | ||
| /// @throws std::runtime_error if the arguments weren't actually valid for the | ||
| /// operation. This should be caught and reported by the testing framework. | ||
| /// @return std::tuple<bool, float> - whether the problem is supported and | ||
| /// kernel execution time (0.0f if s_conf time_kernel is false). | ||
| /// | ||
| /// @see run() | ||
| template <auto SIGNATURE> | ||
| requires ValidConvSignature<SIGNATURE> && ConvDirectionIsForward<SIGNATURE> | ||
| std::tuple<bool, float> run(CkTileConvInstance<SIGNATURE> auto& conv, | ||
| const Args<SIGNATURE>& args, | ||
| const Inputs<SIGNATURE>& inputs, | ||
| const Outputs<SIGNATURE>& outputs, | ||
| const ck_tile::stream_config s_conf = {}) | ||
| { | ||
| using Conv = std::remove_reference_t<decltype(conv)>; | ||
| const auto param = args.to_ck_tile_conv_param(); | ||
|
|
||
| ck_tile::GroupedConvFwdHostArgs<> host_args( | ||
| param, inputs.input, inputs.weight, {}, outputs.output, args.k_batch); | ||
|
|
||
| auto kargs = Conv::MakeKernelArgs(host_args); | ||
|
|
||
| const dim3 grids = Conv::GridSize(kargs); | ||
| const dim3 blocks = Conv::BlockSize(); | ||
|
|
||
| if(!Conv::IsSupportedArgument(kargs)) | ||
| { | ||
| std::cout << "Not supported!"; | ||
| return std::make_tuple(false, 0.f); | ||
| } | ||
|
|
||
| constexpr index_t minimum_occupancy = | ||
| Conv::GemmPipeline::Scheduler == ck_tile::GemmPipelineScheduler::Intrawave ? 1 : 2; | ||
|
|
||
| return std::make_tuple( | ||
| true, | ||
| ck_tile::launch_kernel( | ||
| s_conf, ck_tile::make_kernel<minimum_occupancy>(conv, grids, blocks, 0, kargs))); | ||
| } | ||
|
|
||
| } // namespace ck_tile::builder::test | ||
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.
This lambda looks fine, but isn't this defined some common place? I've seen this repeated other places in the code.
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.
We dont have, I wanted to add tihs but the both usage of this function return different types: vectorck::index_t and vector<ck_tile::index_t>
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.
OK, I'm fine with duplication, we should just think about what were doing each time we duplicate something like this.