diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_dl_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_dl_factory.hpp index 80143427c7..60551b8633 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_dl_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_dl_factory.hpp @@ -25,7 +25,7 @@ struct ConvBwdWeightDlFactory static constexpr size_t SPATIAL_DIM = SIGNATURE.spatial_dim; using Layouts = internal::ConvTensorLayouts; using Types = internal::BwdWeightConvTensorDataTypes; - using Ops = internal::ElementwiseOps; + using Ops = internal::ConvElementwiseOps; using AlgorithmType = decltype(ALGORITHM); static constexpr auto BWD_CONV_SPECIALIZATION = diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_multi_d_wmma_v3_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_multi_d_wmma_v3_factory.hpp index 6f5c679b59..9485e050ef 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_multi_d_wmma_v3_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_multi_d_wmma_v3_factory.hpp @@ -28,7 +28,7 @@ struct ConvBwdWeightMultiDWmmaV3Factory static constexpr size_t SPATIAL_DIM = SIGNATURE.spatial_dim; using Layouts = internal::ConvTensorLayouts; using Types = internal::BwdWeightConvTensorDataTypes; - using Ops = internal::ElementwiseOps; + using Ops = internal::ConvElementwiseOps; using AlgorithmType = decltype(ALGORITHM); static constexpr auto BWD_CONV_SPECIALIZATION = diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_multi_d_xdl_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_multi_d_xdl_factory.hpp index 9f76568ca8..a8c92a39fb 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_multi_d_xdl_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_multi_d_xdl_factory.hpp @@ -28,7 +28,7 @@ struct ConvBwdWeightMultiDXdlFactory static constexpr size_t SPATIAL_DIM = SIGNATURE.spatial_dim; using Layouts = internal::ConvTensorLayouts; using Types = internal::BwdWeightConvTensorDataTypes; - using Ops = internal::ElementwiseOps; + using Ops = internal::ConvElementwiseOps; using AlgorithmType = decltype(ALGORITHM); static constexpr auto BWD_CONV_SPECIALIZATION = diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_two_stage_wmma_v3_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_two_stage_wmma_v3_factory.hpp index 86c48fe322..4f51ac5215 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_two_stage_wmma_v3_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_two_stage_wmma_v3_factory.hpp @@ -28,7 +28,7 @@ struct ConvBwdWeightTwoStageWmmaV3Factory static constexpr size_t SPATIAL_DIM = SIGNATURE.spatial_dim; using Layouts = internal::ConvTensorLayouts; using Types = internal::BwdWeightConvTensorDataTypes; - using Ops = internal::ElementwiseOps; + using Ops = internal::ConvElementwiseOps; using AlgorithmType = decltype(ALGORITHM); static constexpr auto BWD_CONV_SPECIALIZATION = diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_two_stage_xdl_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_two_stage_xdl_factory.hpp index 9c37beae46..3fb04ffa2e 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_two_stage_xdl_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_two_stage_xdl_factory.hpp @@ -28,7 +28,7 @@ struct ConvBwdWeightTwoStageXdlFactory static constexpr size_t SPATIAL_DIM = SIGNATURE.spatial_dim; using Layouts = internal::ConvTensorLayouts; using Types = internal::BwdWeightConvTensorDataTypes; - using Ops = internal::ElementwiseOps; + using Ops = internal::ConvElementwiseOps; using AlgorithmType = decltype(ALGORITHM); static constexpr auto BWD_CONV_SPECIALIZATION = diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_wmma_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_wmma_factory.hpp index 32161a234a..eecc2ad102 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_wmma_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_wmma_factory.hpp @@ -28,7 +28,7 @@ struct ConvBwdWeightWmmaFactory static constexpr size_t SPATIAL_DIM = SIGNATURE.spatial_dim; using Layouts = internal::ConvTensorLayouts; using Types = internal::BwdWeightConvTensorDataTypes; - using Ops = internal::ElementwiseOps; + using Ops = internal::ConvElementwiseOps; using AlgorithmType = decltype(ALGORITHM); static constexpr auto BWD_CONV_SPECIALIZATION = diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_wmma_v3_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_wmma_v3_factory.hpp index baf84402c3..b7845d7a00 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_wmma_v3_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_wmma_v3_factory.hpp @@ -28,7 +28,7 @@ struct ConvBwdWeightWmmaV3Factory static constexpr size_t SPATIAL_DIM = SIGNATURE.spatial_dim; using Layouts = internal::ConvTensorLayouts; using Types = internal::BwdWeightConvTensorDataTypes; - using Ops = internal::ElementwiseOps; + using Ops = internal::ConvElementwiseOps; using AlgorithmType = decltype(ALGORITHM); static constexpr auto BWD_CONV_SPECIALIZATION = diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_xdl_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_xdl_factory.hpp index 91c19d2bd0..db666ffb92 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_xdl_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_xdl_factory.hpp @@ -28,7 +28,7 @@ struct ConvBwdWeightXdlFactory static constexpr size_t SPATIAL_DIM = SIGNATURE.spatial_dim; using Layouts = internal::ConvTensorLayouts; using Types = internal::BwdWeightConvTensorDataTypes; - using Ops = internal::ElementwiseOps; + using Ops = internal::ConvElementwiseOps; using AlgorithmType = decltype(ALGORITHM); static constexpr auto BWD_CONV_SPECIALIZATION = diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_xdl_v3_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_xdl_v3_factory.hpp index f3edd0e6d9..9054aa5b88 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_xdl_v3_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_xdl_v3_factory.hpp @@ -28,7 +28,7 @@ struct ConvBwdWeightXdlV3Factory static constexpr size_t SPATIAL_DIM = SIGNATURE.spatial_dim; using Layouts = internal::ConvTensorLayouts; using Types = internal::BwdWeightConvTensorDataTypes; - using Ops = internal::ElementwiseOps; + using Ops = internal::ConvElementwiseOps; using AlgorithmType = decltype(ALGORITHM); static constexpr auto BWD_CONV_SPECIALIZATION = diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_fwd_dl_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_fwd_dl_factory.hpp index 31246eb5a8..679ce4e59e 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_fwd_dl_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_fwd_dl_factory.hpp @@ -26,7 +26,7 @@ struct ConvFwdDlFactory static constexpr size_t SPATIAL_DIM = SIGNATURE.spatial_dim; using Layouts = internal::ConvTensorLayouts; using Types = internal::FwdConvTensorDataTypes; - using Ops = internal::ElementwiseOps; + using Ops = internal::ConvElementwiseOps; using AlgorithmType = decltype(ALGORITHM); static constexpr auto FWD_CONV_SPECIALIZATION = internal::SetFwdConvSpecialization(); @@ -94,13 +94,13 @@ struct ConvFwdDlFactory typename Types::DsDataTypes, typename Types::EDataType, typename Types::AccDataType, - typename Layouts::ALayout, - typename Layouts::BLayout, + typename Layouts::InLayout, + typename Layouts::WeiLayout, typename Layouts::DsLayout, - typename Layouts::ELayout, - typename Ops::AElementwiseOp, - typename Ops::BElementwiseOp, - typename Ops::CDEElementwiseOp, + typename Layouts::OutLayout, + typename Ops::InElementwiseOp, + typename Ops::WeiElementwiseOp, + typename Ops::OutElementwiseOp, FWD_CONV_SPECIALIZATION, GEMM_SPECIALIZATION, BLOCK.block_size, diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_fwd_large_tensor_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_fwd_large_tensor_factory.hpp index 9cd56ad7ad..1ad00d5c7a 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_fwd_large_tensor_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_fwd_large_tensor_factory.hpp @@ -28,7 +28,7 @@ struct ConvFwdLargeTensorFactory static constexpr size_t SPATIAL_DIM = SIGNATURE.spatial_dim; using Layouts = internal::ConvTensorLayouts; using Types = internal::FwdConvTensorDataTypes; - using Ops = internal::ElementwiseOps; + using Ops = internal::ConvElementwiseOps; using AlgorithmType = decltype(ALGORITHM); static constexpr auto FWD_CONV_SPECIALIZATION = internal::SetFwdConvSpecialization(); @@ -59,19 +59,19 @@ struct ConvFwdLargeTensorFactory using Instance = ck::tensor_operation::device::DeviceGroupedConvFwdMultipleD_Xdl_CShuffle_Large_Tensor< SPATIAL_DIM, - typename Layouts::ALayout, - typename Layouts::BLayout, + typename Layouts::InLayout, + typename Layouts::WeiLayout, typename Layouts::DsLayout, - typename Layouts::ELayout, + typename Layouts::OutLayout, typename Types::ADataType, typename Types::BDataType, typename Types::AccDataType, typename Types::CShuffleDataType, typename Types::DsDataTypes, typename Types::EDataType, - typename Ops::AElementwiseOp, - typename Ops::BElementwiseOp, - typename Ops::CDEElementwiseOp, + typename Ops::InElementwiseOp, + typename Ops::WeiElementwiseOp, + typename Ops::OutElementwiseOp, SPECIALIZATION.conv_spec, SPECIALIZATION.gemm_spec, ALGORITHM.num_gemm_k_prefetch_stages, diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_fwd_v3_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_fwd_v3_factory.hpp index 7d889a0c01..210f439316 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_fwd_v3_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_fwd_v3_factory.hpp @@ -28,7 +28,7 @@ struct ConvFwdXdlV3Factory static constexpr size_t SPATIAL_DIM = SIGNATURE.spatial_dim; using Layouts = internal::ConvTensorLayouts; using Types = internal::FwdConvTensorDataTypes; - using Ops = internal::ElementwiseOps; + using Ops = internal::ConvElementwiseOps; using AlgorithmType = decltype(ALGORITHM); static_assert(ALGORITHM.transfer.a.lds_transfer.is_direct_load == @@ -64,19 +64,19 @@ struct ConvFwdXdlV3Factory // The forward convolution kernel class instance. using Instance = ck::tensor_operation::device::DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3< SPATIAL_DIM, - typename Layouts::ALayout, - typename Layouts::BLayout, + typename Layouts::InLayout, + typename Layouts::WeiLayout, typename Layouts::DsLayout, - typename Layouts::ELayout, + typename Layouts::OutLayout, typename Types::ADataType, typename Types::BDataType, typename Types::AccDataType, typename Types::CShuffleDataType, typename Types::DsDataTypes, typename Types::EDataType, - typename Ops::AElementwiseOp, - typename Ops::BElementwiseOp, - typename Ops::CDEElementwiseOp, + typename Ops::InElementwiseOp, + typename Ops::WeiElementwiseOp, + typename Ops::OutElementwiseOp, SPECIALIZATION.conv_spec, SPECIALIZATION.gemm_spec, BLOCK.block_size, diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_fwd_wmma_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_fwd_wmma_factory.hpp index 3506f5d1a9..01a78738ce 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_fwd_wmma_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_fwd_wmma_factory.hpp @@ -28,7 +28,7 @@ struct ConvFwdWmmaFactory static constexpr size_t SPATIAL_DIM = SIGNATURE.spatial_dim; using Layouts = internal::ConvTensorLayouts; using Types = internal::FwdConvTensorDataTypes; - using Ops = internal::ElementwiseOps; + using Ops = internal::ConvElementwiseOps; using AlgorithmType = decltype(ALGORITHM); static constexpr auto FWD_CONV_SPECIALIZATION = internal::SetFwdConvSpecialization(); @@ -60,19 +60,19 @@ struct ConvFwdWmmaFactory // The forward convolution kernel class instance. using Instance = ck::tensor_operation::device::DeviceGroupedConvFwdMultipleD_Wmma_CShuffle< SPATIAL_DIM, - typename Layouts::ALayout, - typename Layouts::BLayout, + typename Layouts::InLayout, + typename Layouts::WeiLayout, typename Layouts::DsLayout, - typename Layouts::ELayout, + typename Layouts::OutLayout, typename Types::ADataType, typename Types::BDataType, typename Types::AccDataType, typename Types::CShuffleDataType, typename Types::DsDataTypes, typename Types::EDataType, - typename Ops::AElementwiseOp, - typename Ops::BElementwiseOp, - typename Ops::CDEElementwiseOp, + typename Ops::InElementwiseOp, + typename Ops::WeiElementwiseOp, + typename Ops::OutElementwiseOp, SPECIALIZATION.conv_spec, SPECIALIZATION.gemm_spec, ALGORITHM.num_gemm_k_prefetch_stages, diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_fwd_xdl_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_fwd_xdl_factory.hpp index 446ceceda2..50116a4f87 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_fwd_xdl_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_fwd_xdl_factory.hpp @@ -28,7 +28,7 @@ struct ConvFwdXdlFactory static constexpr size_t SPATIAL_DIM = SIGNATURE.spatial_dim; using Layouts = internal::ConvTensorLayouts; using Types = internal::FwdConvTensorDataTypes; - using Ops = internal::ElementwiseOps; + using Ops = internal::ConvElementwiseOps; using AlgorithmType = decltype(ALGORITHM); static constexpr auto FWD_CONV_SPECIALIZATION = internal::SetFwdConvSpecialization(); @@ -59,19 +59,19 @@ struct ConvFwdXdlFactory // The forward convolution kernel class instance. using Instance = ck::tensor_operation::device::DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle< SPATIAL_DIM, - typename Layouts::ALayout, - typename Layouts::BLayout, + typename Layouts::InLayout, + typename Layouts::WeiLayout, typename Layouts::DsLayout, - typename Layouts::ELayout, + typename Layouts::OutLayout, typename Types::ADataType, typename Types::BDataType, typename Types::AccDataType, typename Types::CShuffleDataType, typename Types::DsDataTypes, typename Types::EDataType, - typename Ops::AElementwiseOp, - typename Ops::BElementwiseOp, - typename Ops::CDEElementwiseOp, + typename Ops::InElementwiseOp, + typename Ops::WeiElementwiseOp, + typename Ops::OutElementwiseOp, SPECIALIZATION.conv_spec, SPECIALIZATION.gemm_spec, ALGORITHM.num_gemm_k_prefetch_stages, diff --git a/experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_elementwise_op.hpp b/experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_elementwise_op.hpp index c7d9ce0ac6..d2f82d3ecd 100644 --- a/experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_elementwise_op.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_elementwise_op.hpp @@ -62,30 +62,20 @@ consteval auto GetElementwiseOp() } template -struct ElementwiseOps +struct ConvElementwiseOps { - private: static constexpr auto input_op = GetElementwiseOp(); static constexpr auto weight_op = GetElementwiseOp(); static constexpr auto output_op = GetElementwiseOp(); - static constexpr bool is_forward = ConvDirectionIsForward; - static constexpr bool is_bwd_weight = ConvDirectionIsBackwardWeight; - - using InputOp = typename decltype(input_op)::Op; - using WeightOp = typename decltype(weight_op)::Op; - using OutputOp = typename decltype(output_op)::Op; - - public: - // Forward convolution elementwise ops - using AElementwiseOp = std::conditional_t; - using BElementwiseOp = std::conditional_t; - using CDEElementwiseOp = std::conditional_t; - - // Backward weight convolution elementwise ops - using InElementwiseOp = std::conditional_t; - using WeiElementwiseOp = std::conditional_t; - using OutElementwiseOp = std::conditional_t; + using InElementwiseOp = typename decltype(input_op)::Op; + using WeiElementwiseOp = typename decltype(weight_op)::Op; + using OutElementwiseOp = typename decltype(output_op)::Op; + + // TODO: Remove, now left for compatibility. Factories do not need it anymore. + // using AElementwiseOp = InElementwiseOp; + // using BElementwiseOp = WeiElementwiseOp; + // using CDEElementwiseOp = OutElementwiseOp; }; } // namespace ck_tile::builder::factory::internal diff --git a/experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_tensor_layout.hpp b/experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_tensor_layout.hpp index d08dddff83..f44ce15494 100644 --- a/experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_tensor_layout.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_tensor_layout.hpp @@ -222,28 +222,15 @@ template ValidConvOutputLayoutForSpatialDim) struct ConvTensorLayouts { - private: - static constexpr bool is_forward = ConvDirectionIsForward; - static constexpr bool is_bwd_weight = ConvDirectionIsBackwardWeight; + using InLayout = decltype(TensorLayoutToCK()); + using WeiLayout = decltype(TensorLayoutToCK()); + using OutLayout = decltype(TensorLayoutToCK()); + using DsLayout = decltype(GetAuxiliaryTensorLayouts())::type; - using InputLayout = decltype(TensorLayoutToCK()); - using WeightLayout = decltype(TensorLayoutToCK()); - using OutputLayout = decltype(TensorLayoutToCK()); - using AuxLayout = decltype(GetAuxiliaryTensorLayouts())::type; - - public: - // Forward convolution layouts - using ALayout = std::conditional_t; - using BLayout = std::conditional_t; - using ELayout = std::conditional_t; - - // Backward weight convolution layouts - using InLayout = std::conditional_t; - using WeiLayout = std::conditional_t; - using OutLayout = std::conditional_t; - - // Applicable for all directions - using DsLayout = AuxLayout; + // TODO: Remove,now left for compatibility. Factories do not need it anymore. + // using ALayout = InLayout; + // using BLayout = WeiLayout; + // using ELayout = OutLayout; }; } // namespace ck_tile::builder::factory::internal diff --git a/experimental/builder/include/ck_tile/builder/testing/conv_fwd.hpp b/experimental/builder/include/ck_tile/builder/testing/conv_fwd.hpp index 3240033c55..28024cd1b8 100644 --- a/experimental/builder/include/ck_tile/builder/testing/conv_fwd.hpp +++ b/experimental/builder/include/ck_tile/builder/testing/conv_fwd.hpp @@ -71,7 +71,7 @@ struct Args using OutputDescriptor = TensorDescriptor; // TODO: We shouldn't need to call into an internal namespace here. - using Ops = factory::internal::ElementwiseOps; + using Ops = factory::internal::ConvElementwiseOps; // TODO: We shouldn't need to call into an internal namespace here. using Layouts = factory::internal::ConvTensorLayouts; @@ -88,9 +88,9 @@ struct Args FilterExtent input_left_pad; FilterExtent input_right_pad; - Ops::AElementwiseOp a_elementwise_op; - Ops::BElementwiseOp b_elementwise_op; - Ops::CDEElementwiseOp cde_elementwise_op; + Ops::InElementwiseOp a_elementwise_op; + Ops::WeiElementwiseOp b_elementwise_op; + Ops::OutElementwiseOp cde_elementwise_op; /// This function returns the `TensorDescriptor` corresponding to /// the input-tensor of the convolution problem. This can then @@ -105,7 +105,7 @@ struct Args // function. const auto param = to_ck_conv_param(); const auto desc = ck::utils::conv::make_input_host_tensor_descriptor_g_n_c_wis_packed< - typename Layouts::ALayout>(param); + typename Layouts::InLayout>(param); using Extent = typename InputDescriptor::Extent; return InputDescriptor(Extent::from_vector(desc.GetLengths()), Extent::from_vector(desc.GetStrides())); @@ -119,7 +119,7 @@ struct Args // See note in implementation of `make_input_descriptor`. const auto param = to_ck_conv_param(); const auto desc = ck::utils::conv::make_weight_host_tensor_descriptor_g_k_c_xs_packed< - typename Layouts::BLayout>(param); + typename Layouts::WeiLayout>(param); using Extent = typename WeightDescriptor::Extent; return WeightDescriptor(Extent::from_vector(desc.GetLengths()), Extent::from_vector(desc.GetStrides())); @@ -133,7 +133,7 @@ struct Args // See note in implementation of `make_input_descriptor`. const auto param = to_ck_conv_param(); const auto desc = ck::utils::conv::make_output_host_tensor_descriptor_g_n_k_wos_packed< - typename Layouts::ELayout>(param); + typename Layouts::OutLayout>(param); using Extent = typename OutputDescriptor::Extent; return OutputDescriptor(Extent::from_vector(desc.GetLengths()), Extent::from_vector(desc.GetStrides())); diff --git a/experimental/builder/include/ck_tile/builder/testing/conv_fwd_ck.hpp b/experimental/builder/include/ck_tile/builder/testing/conv_fwd_ck.hpp index 499e0ef3de..a90f53ba7d 100644 --- a/experimental/builder/include/ck_tile/builder/testing/conv_fwd_ck.hpp +++ b/experimental/builder/include/ck_tile/builder/testing/conv_fwd_ck.hpp @@ -27,7 +27,7 @@ template > + typename Ops = factory::internal::ConvElementwiseOps> concept CkConvInstance = requires(Conv& conv, // TODO: This should be changed depending on IsMultiA etc. // Currently that is not yet supported elsewhere anyway. @@ -37,9 +37,9 @@ concept CkConvInstance = requires(Conv& conv, std::array lengths, std::array strides, std::array filter, - Ops::AElementwiseOp elementwise_a, - Ops::BElementwiseOp elementwise_b, - Ops::CDEElementwiseOp elementwise_cde) { + Ops::InElementwiseOp elementwise_a, + Ops::WeiElementwiseOp elementwise_b, + Ops::OutElementwiseOp elementwise_cde) { { conv.MakeArgument(p_a, p_b, diff --git a/experimental/builder/test/unit_conv_tensor_layout.cpp b/experimental/builder/test/unit_conv_tensor_layout.cpp index 8c1ba5562e..0df94d977e 100644 --- a/experimental/builder/test/unit_conv_tensor_layout.cpp +++ b/experimental/builder/test/unit_conv_tensor_layout.cpp @@ -40,9 +40,9 @@ TEST(ConvTensorLayout, AssignsLayoutsFor1D_NWGC_GKXC_NWGK) using TensorLayouts = ConvTensorLayouts; - EXPECT_TRUE((std::is_same_v)); - EXPECT_TRUE((std::is_same_v)); - EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v>)); } @@ -59,9 +59,9 @@ TEST(ConvTensorLayout, AssignsLayoutsFor1D_NGCW_GKXC_NGKW) using TensorLayouts = ConvTensorLayouts; - EXPECT_TRUE((std::is_same_v)); - EXPECT_TRUE((std::is_same_v)); - EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v>)); } @@ -78,9 +78,9 @@ TEST(ConvTensorLayout, AssignsLayoutsFor1D_GNWC_GKXC_GNWK) using TensorLayouts = ConvTensorLayouts; - EXPECT_TRUE((std::is_same_v)); - EXPECT_TRUE((std::is_same_v)); - EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v>)); } @@ -97,9 +97,9 @@ TEST(ConvTensorLayout, AssignsLayoutsFor1D_NGCW_GKCX_NGKW) using TensorLayouts = ConvTensorLayouts; - EXPECT_TRUE((std::is_same_v)); - EXPECT_TRUE((std::is_same_v)); - EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v>)); } @@ -116,9 +116,9 @@ TEST(ConvTensorLayout, AssignsLayoutsFor2D_NGCHW_GKYXC_NGKHW) using TensorLayouts = ConvTensorLayouts; - EXPECT_TRUE((std::is_same_v)); - EXPECT_TRUE((std::is_same_v)); - EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v>)); } @@ -135,9 +135,9 @@ TEST(ConvTensorLayout, AssignsLayoutsFor2D_NHWGC_GKYXC_NHWGK) using TensorLayouts = ConvTensorLayouts; - EXPECT_TRUE((std::is_same_v)); - EXPECT_TRUE((std::is_same_v)); - EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v>)); } @@ -154,9 +154,9 @@ TEST(ConvTensorLayout, AssignsLayoutsFor2D_GNHWC_GKYXC_GNHWK) using TensorLayouts = ConvTensorLayouts; - EXPECT_TRUE((std::is_same_v)); - EXPECT_TRUE((std::is_same_v)); - EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v>)); } @@ -173,9 +173,9 @@ TEST(ConvTensorLayout, AssignsLayoutsFor2D_NGCHW_GKCYX_NGKHW) using TensorLayouts = ConvTensorLayouts; - EXPECT_TRUE((std::is_same_v)); - EXPECT_TRUE((std::is_same_v)); - EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v>)); } @@ -192,9 +192,9 @@ TEST(ConvTensorLayout, AssignsLayoutsFor3D_NGCDHW_GKCZYX_NGKDHW) using TensorLayouts = ConvTensorLayouts; - EXPECT_TRUE((std::is_same_v)); - EXPECT_TRUE((std::is_same_v)); - EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v>)); } @@ -211,9 +211,9 @@ TEST(ConvTensorLayout, AssignsLayoutsFor3D_NDHWGC_GKZYXC_NDHWGK) using TensorLayouts = ConvTensorLayouts; - EXPECT_TRUE((std::is_same_v)); - EXPECT_TRUE((std::is_same_v)); - EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v>)); } @@ -230,9 +230,9 @@ TEST(ConvTensorLayout, AssignsLayoutsFor3D_GNDHWC_GKZYXC_GNDHWK) using TensorLayouts = ConvTensorLayouts; - EXPECT_TRUE((std::is_same_v)); - EXPECT_TRUE((std::is_same_v)); - EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v>)); } @@ -389,9 +389,9 @@ TEST(ConvTensorLayoutsWithAuxiliary, Conv2DWithSingleBiasG_K) using TensorLayouts = ConvTensorLayouts; - EXPECT_TRUE((std::is_same_v)); - EXPECT_TRUE((std::is_same_v)); - EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); using ExpectedDsLayout = ck::Tuple; EXPECT_TRUE((std::is_same_v)); @@ -416,9 +416,9 @@ TEST(ConvTensorLayoutsWithAuxiliary, Conv2DWithSingleBiasGC) using TensorLayouts = ConvTensorLayouts; - EXPECT_TRUE((std::is_same_v)); - EXPECT_TRUE((std::is_same_v)); - EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); using ExpectedDsLayout = ck::Tuple; EXPECT_TRUE((std::is_same_v)); @@ -444,9 +444,9 @@ TEST(ConvTensorLayoutsWithAuxiliary, Conv2DWithTwoAuxiliaryTensors) using TensorLayouts = ConvTensorLayouts; - EXPECT_TRUE((std::is_same_v)); - EXPECT_TRUE((std::is_same_v)); - EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); using ExpectedDsLayout = ck::Tuple; @@ -472,9 +472,9 @@ TEST(ConvTensorLayoutsWithAuxiliary, Conv1DWithBias) using TensorLayouts = ConvTensorLayouts; - EXPECT_TRUE((std::is_same_v)); - EXPECT_TRUE((std::is_same_v)); - EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); using ExpectedDsLayout = ck::Tuple; EXPECT_TRUE((std::is_same_v)); @@ -499,9 +499,9 @@ TEST(ConvTensorLayoutsWithAuxiliary, Conv3DWithBias) using TensorLayouts = ConvTensorLayouts; - EXPECT_TRUE((std::is_same_v)); - EXPECT_TRUE((std::is_same_v)); - EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); + EXPECT_TRUE((std::is_same_v)); using ExpectedDsLayout = ck::Tuple; EXPECT_TRUE((std::is_same_v));