ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-02 21:27:45 +00:00
Code Issues Packages Projects Releases Wiki Activity

Value constraint concept example.

Browse Source
This commit is contained in:
Ville Pietilä
2025-10-15 07:46:47 +00:00
parent 12c812e993
commit 28221cf01f
2 changed files with 14 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files

11
experimental/builder/include/ck_tile/builder/conv_algorithm.hpp
View File

@@ -86,12 +86,19 @@ concept VectorTransferDescriptorAB = requires(T t) {
};
// Concept for the C tensor vectors transfer details.
template <auto Value>
concept ValidVectorTransferC = requires {
requires Value.scaler_per_vector > 0 &&
Value.m_xdl_per_wave_per_shuffle > 0 &&
Value.n_xdl_per_wave_per_shuffle > 0 ;
};
template <typename T>
concept VectorTransferDescriptorC = requires(T t) {
{ t.m_xdl_per_wave_per_shuffle } -> std::convertible_to<size_t>;
{ t.n_xdl_per_wave_per_shuffle } -> std::convertible_to<size_t>;
{ t.scaler_per_vector } -> std::convertible_to<size_t>;
};
{ t.scaler_per_vector } -> std::convertible_to<size_t>;
};
// Concept to check if a struct specifies A Block tranfer info.
template <typename T>

5
experimental/builder/include/ck_tile/builder/conv_factory.hpp
View File

@@ -33,6 +33,7 @@
#include <ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_abd_xdl_cshuffle_v3.hpp>
#include <ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_data_multiple_d_xdl_cshuffle_v1.hpp>
#include <ck_tile/builder/conv_signature.hpp>
#include <ck_tile/builder/conv_algorithm.hpp>
#include <ck_tile/builder/conv_algorithm_types.hpp>
#include <ck_tile/builder/builder_utils.hpp>
#include <ck_tile/builder/types.hpp>
@@ -406,6 +407,10 @@ struct ConvFactory<SIGNATURE, ALGORITHM, VERSION>
static constexpr auto C_BLOCK_TRANSFER = factory_internal::SetCBlockTransfer<SIGNATURE, ALGORITHM>();
static constexpr auto PIPELINE_SCHEDULER = ck::BlockGemmPipelineScheduler::Intrawave;
static constexpr auto PIPELINE_VERSION = factory_internal::SetBlockGemmPipelineVersion<ALGORITHM>();
// Preconditions
static_assert(ValidVectorTransferC<C_BLOCK_TRANSFER>);
// The forward convolution kernel class instance.
using Instance =
ck::tensor_operation::device::DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3< //