Merge branch 'develop' into aosewski/test_ggemm_splitk

client_example/18_groupnorm/groupnorm_swish.cpp

@@ -131,11 +131,12 @@ int main(int argc, char* argv[])
         }
     }
 
-    std::cout << "Best Perf: " << best_ave_time << " ms, " << best_gb_per_sec << " GB/s, "
-              << best_op_name << std::endl;
-
     // run the best intance
+    if(found)
     {
+        std::cout << "Best Perf: " << best_ave_time << " ms, " << best_gb_per_sec << " GB/s, "
+                  << best_op_name << std::endl;
+
         auto& op_ptr = op_ptrs[best_op_id];
         std::cout << "Run the best instance without timing: " << op_ptr->GetTypeString()
                   << std::endl;

client_example/19_pool_fwd/CMakeLists.txt

+add_executable(client_max_pool2d_fwd max_pool2d_fwd.cpp)
+target_link_libraries(client_max_pool2d_fwd PRIVATE composable_kernel::device_operations)
+
+add_executable(client_avg_pool3d_fwd avg_pool3d_fwd.cpp)
+target_link_libraries(client_avg_pool3d_fwd PRIVATE composable_kernel::device_operations)
\ No newline at end of file

client_example/19_pool_fwd/avg_pool3d_fwd.cpp

+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
+
+#include <iomanip>
+#include <vector>
+#include <iostream>
+
+#include "ck/ck.hpp"
+#include "ck/tensor_operation/gpu/device/tensor_layout.hpp"
+#include "ck/tensor_operation/gpu/device/device_pool_fwd.hpp"
+#include "ck/tensor_operation/gpu/element/element_wise_operation.hpp"
+
+#include "ck/library/tensor_operation_instance/gpu/pool3d_fwd.hpp"
+
+using InDataType    = ck::half_t;
+using OutDataType   = ck::half_t;
+using IndexDataType = int32_t;
+
+constexpr ck::index_t InOutRank  = 5;
+constexpr ck::index_t WindowRank = 3;
+#if 0
+constexpr auto ReduceOpId  = ck::ReduceTensorOp::MAX;
+constexpr bool OutputIndex = false;
+#else
+constexpr auto ReduceOpId  = ck::ReduceTensorOp::AVG;
+constexpr bool OutputIndex = false;
+#endif
+
+struct SimpleDeviceMem
+{
+    SimpleDeviceMem() = delete;
+
+    SimpleDeviceMem(std::size_t mem_size) : p_mem_{}
+    {
+        (void)hipMalloc(static_cast<void**>(&p_mem_), mem_size);
+    }
+
+    void* GetDeviceBuffer() { return p_mem_; }
+
+    ~SimpleDeviceMem() { (void)hipFree(p_mem_); }
+
+    void* p_mem_;
+};
+
+int main(int argc, char* argv[])
+{
+    ck::index_t N               = 2;
+    ck::index_t C               = 32;
+    ck::index_t Z               = 2;
+    ck::index_t Y               = 2;
+    ck::index_t X               = 2;
+    ck::index_t Di              = 30;
+    ck::index_t Hi              = 30;
+    ck::index_t Wi              = 30;
+    ck::index_t window_stride_d = 2;
+    ck::index_t window_stride_h = 2;
+    ck::index_t window_stride_w = 2;
+    ck::index_t in_left_pad_d   = 1;
+    ck::index_t in_left_pad_h   = 1;
+    ck::index_t in_left_pad_w   = 1;
+    ck::index_t in_right_pad_d  = 1;
+    ck::index_t in_right_pad_h  = 1;
+    ck::index_t in_right_pad_w  = 1;
+
+    ck::index_t Do = (Di + in_left_pad_d + in_right_pad_d - Z) / window_stride_d + 1;
+    ck::index_t Ho = (Hi + in_left_pad_h + in_right_pad_h - Y) / window_stride_h + 1;
+    ck::index_t Wo = (Wi + in_left_pad_w + in_right_pad_w - X) / window_stride_w + 1;
+
+    // Pool API only support the order of NCDHW
+    std::vector<ck::index_t> in_length              = {N, C, Di, Hi, Wi};
+    std::vector<ck::index_t> out_length             = {N, C, Do, Ho, Wo};
+    std::vector<ck::index_t> window_spatial_lengths = {Z, Y, X};
+    std::vector<ck::index_t> window_strides   = {window_stride_d, window_stride_h, window_stride_w};
+    std::vector<ck::index_t> input_left_pads  = {in_left_pad_d, in_left_pad_h, in_left_pad_w};
+    std::vector<ck::index_t> input_right_pads = {in_right_pad_d, in_right_pad_h, in_right_pad_w};
+
+    std::size_t in_tensor_size  = N * C * Di * Hi * Wi;
+    std::size_t out_tensor_size = N * C * Do * Ho * Wo;
+
+    // tensor layout = NDHWC
+    std::vector<ck::index_t> in_tensor_stride  = {Di * C * Hi * Wi, 1, C * Hi * Wi, Wi * C, C};
+    std::vector<ck::index_t> out_tensor_stride = {Do * C * Ho * Wo, 1, C * Ho * Wo, Wo * C, C};
+
+    SimpleDeviceMem in_device_buf(sizeof(InDataType) * in_tensor_size);
+    SimpleDeviceMem out_device_buf(sizeof(OutDataType) * out_tensor_size);
+    SimpleDeviceMem out_indices_device_buf(sizeof(IndexDataType) * out_tensor_size);
+
+    using DeviceOp = ck::tensor_operation::device::DevicePoolFwd<InOutRank,
+                                                                 WindowRank,
+                                                                 InDataType,
+                                                                 OutDataType,
+                                                                 IndexDataType,
+                                                                 ReduceOpId,
+                                                                 OutputIndex>;
+
+    // get device op instances
+    const auto op_ptrs = ck::tensor_operation::device::instance::DeviceOperationInstanceFactory<
+        DeviceOp>::GetInstances();
+
+    std::cout << "found " << op_ptrs.size() << " instances" << std::endl;
+
+    std::string best_op_name;
+    bool found            = false;
+    int best_op_id        = -1;
+    float best_ave_time   = std::numeric_limits<float>::max();
+    float best_gb_per_sec = 0;
+
+    // profile device operation instances
+    std::cout << "Run all instances and do timing" << std::endl;
+
+    for(int i = 0; i < op_ptrs.size(); ++i)
+    {
+        auto& op_ptr      = op_ptrs[i];
+        auto argument_ptr = op_ptr->MakeArgumentPointer(
+            static_cast<InDataType*>(in_device_buf.GetDeviceBuffer()),
+            static_cast<OutDataType*>(out_device_buf.GetDeviceBuffer()),
+            static_cast<IndexDataType*>(out_indices_device_buf.GetDeviceBuffer()),
+            in_length,
+            window_spatial_lengths,
+            out_length,
+            in_tensor_stride,
+            out_tensor_stride,
+            out_tensor_stride,
+            window_strides,
+            input_left_pads,
+            input_right_pads,
+            {2, 3, 4});
+
+        auto invoker_ptr = op_ptr->MakeInvokerPointer();
+
+        std::string op_name = op_ptr->GetTypeString();
+
+        if(op_ptr->IsSupportedArgument(argument_ptr.get()))
+        {
+            float ave_time = invoker_ptr->Run(argument_ptr.get(), StreamConfig{nullptr, true});
+
+            std::size_t num_bytes =
+                in_tensor_size * sizeof(InDataType) + out_tensor_size * sizeof(OutDataType);
+
+            if constexpr(OutputIndex)
+                num_bytes += out_tensor_size * sizeof(IndexDataType);
+
+            float gb_per_sec = num_bytes / 1.E6 / ave_time;
+
+            std::cout << "Perf: " << std::setw(10) << ave_time << " ms, " << gb_per_sec << " GB/s, "
+                      << op_name << std::endl;
+
+            if(ave_time < best_ave_time)
+            {
+                found           = true;
+                best_op_id      = i;
+                best_op_name    = op_name;
+                best_ave_time   = ave_time;
+                best_gb_per_sec = gb_per_sec;
+            }
+        }
+        else
+        {
+            std::cout << op_name << " does not support this problem" << std::endl;
+        }
+    }
+
+    // run the best intance
+    if(found)
+    {
+        std::cout << "Best Perf: " << best_ave_time << " ms, " << best_gb_per_sec << " GB/s, "
+                  << best_op_name << std::endl;
+
+        auto& op_ptr = op_ptrs[best_op_id];
+        std::cout << "Run the best instance without timing: " << op_ptr->GetTypeString()
+                  << std::endl;
+
+        auto argument_ptr = op_ptr->MakeArgumentPointer(
+            static_cast<InDataType*>(in_device_buf.GetDeviceBuffer()),
+            static_cast<OutDataType*>(out_device_buf.GetDeviceBuffer()),
+            static_cast<IndexDataType*>(out_indices_device_buf.GetDeviceBuffer()),
+            in_length,
+            window_spatial_lengths,
+            out_length,
+            in_tensor_stride,
+            out_tensor_stride,
+            out_tensor_stride,
+            window_strides,
+            input_left_pads,
+            input_right_pads,
+            {2, 3, 4});
+
+        auto invoker_ptr = op_ptr->MakeInvokerPointer();
+
+        if(op_ptr->IsSupportedArgument(argument_ptr.get()))
+        {
+            invoker_ptr->Run(argument_ptr.get(), StreamConfig{nullptr, false});
+        }
+
+        std::cout << "Done" << std::endl;
+    }
+
+    return 0;
+}

client_example/19_pool_fwd/max_pool2d_fwd.cpp

+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
+
+#include <iomanip>
+#include <vector>
+#include <iostream>
+
+#include "ck/ck.hpp"
+#include "ck/tensor_operation/gpu/device/tensor_layout.hpp"
+#include "ck/tensor_operation/gpu/device/device_pool_fwd.hpp"
+#include "ck/tensor_operation/gpu/element/element_wise_operation.hpp"
+
+#include "ck/library/tensor_operation_instance/gpu/pool2d_fwd.hpp"
+
+using InDataType    = ck::half_t;
+using OutDataType   = ck::half_t;
+using IndexDataType = int32_t;
+
+constexpr ck::index_t InOutRank  = 4;
+constexpr ck::index_t WindowRank = 2;
+#if 1
+constexpr auto ReduceOpId  = ck::ReduceTensorOp::MAX;
+constexpr bool OutputIndex = true;
+#else
+constexpr auto ReduceOpId  = ck::ReduceTensorOp::AVG;
+constexpr bool OutputIndex = false;
+#endif
+
+struct SimpleDeviceMem
+{
+    SimpleDeviceMem() = delete;
+
+    SimpleDeviceMem(std::size_t mem_size) : p_mem_{}
+    {
+        (void)hipMalloc(static_cast<void**>(&p_mem_), mem_size);
+    }
+
+    void* GetDeviceBuffer() { return p_mem_; }
+
+    ~SimpleDeviceMem() { (void)hipFree(p_mem_); }
+
+    void* p_mem_;
+};
+
+int main(int argc, char* argv[])
+{
+    ck::index_t N               = 2;
+    ck::index_t C               = 32;
+    ck::index_t Y               = 2;
+    ck::index_t X               = 2;
+    ck::index_t Hi              = 30;
+    ck::index_t Wi              = 30;
+    ck::index_t window_stride_h = 2;
+    ck::index_t window_stride_w = 2;
+    ck::index_t in_left_pad_h   = 1;
+    ck::index_t in_left_pad_w   = 1;
+    ck::index_t in_right_pad_h  = 1;
+    ck::index_t in_right_pad_w  = 1;
+
+    ck::index_t Ho = (Hi + in_left_pad_h + in_right_pad_h - Y) / window_stride_h + 1;
+    ck::index_t Wo = (Wi + in_left_pad_w + in_right_pad_w - X) / window_stride_w + 1;
+
+    // Pool API only support the order of NCHW
+    std::vector<ck::index_t> in_length              = {N, C, Hi, Wi};
+    std::vector<ck::index_t> out_length             = {N, C, Ho, Wo};
+    std::vector<ck::index_t> window_spatial_lengths = {Y, X};
+    std::vector<ck::index_t> window_strides         = {window_stride_h, window_stride_w};
+    std::vector<ck::index_t> input_left_pads        = {in_left_pad_h, in_left_pad_w};
+    std::vector<ck::index_t> input_right_pads       = {in_right_pad_h, in_right_pad_w};
+
+    std::size_t in_tensor_size  = N * C * Hi * Wi;
+    std::size_t out_tensor_size = N * C * Ho * Wo;
+
+    // tensor layout = NHWC
+    std::vector<ck::index_t> in_tensor_stride  = {C * Hi * Wi, 1, Wi * C, C};
+    std::vector<ck::index_t> out_tensor_stride = {C * Ho * Wo, 1, Wo * C, C};
+
+    SimpleDeviceMem in_device_buf(sizeof(InDataType) * in_tensor_size);
+    SimpleDeviceMem out_device_buf(sizeof(OutDataType) * out_tensor_size);
+    SimpleDeviceMem out_indices_device_buf(sizeof(IndexDataType) * out_tensor_size);
+
+    using DeviceOp = ck::tensor_operation::device::DevicePoolFwd<InOutRank,
+                                                                 WindowRank,
+                                                                 InDataType,
+                                                                 OutDataType,
+                                                                 IndexDataType,
+                                                                 ReduceOpId,
+                                                                 OutputIndex>;
+
+    // get device op instances
+    const auto op_ptrs = ck::tensor_operation::device::instance::DeviceOperationInstanceFactory<
+        DeviceOp>::GetInstances();
+
+    std::cout << "found " << op_ptrs.size() << " instances" << std::endl;
+
+    std::string best_op_name;
+    bool found            = false;
+    int best_op_id        = -1;
+    float best_ave_time   = std::numeric_limits<float>::max();
+    float best_gb_per_sec = 0;
+
+    // profile device operation instances
+    std::cout << "Run all instances and do timing" << std::endl;
+
+    for(int i = 0; i < op_ptrs.size(); ++i)
+    {
+        auto& op_ptr      = op_ptrs[i];
+        auto argument_ptr = op_ptr->MakeArgumentPointer(
+            static_cast<InDataType*>(in_device_buf.GetDeviceBuffer()),
+            static_cast<OutDataType*>(out_device_buf.GetDeviceBuffer()),
+            static_cast<IndexDataType*>(out_indices_device_buf.GetDeviceBuffer()),
+            in_length,
+            window_spatial_lengths,
+            out_length,
+            in_tensor_stride,
+            out_tensor_stride,
+            out_tensor_stride,
+            window_strides,
+            input_left_pads,
+            input_right_pads,
+            {2, 3});
+
+        auto invoker_ptr = op_ptr->MakeInvokerPointer();
+
+        std::string op_name = op_ptr->GetTypeString();
+
+        if(op_ptr->IsSupportedArgument(argument_ptr.get()))
+        {
+            float ave_time = invoker_ptr->Run(argument_ptr.get(), StreamConfig{nullptr, true});
+
+            std::size_t num_bytes =
+                in_tensor_size * sizeof(InDataType) + out_tensor_size * sizeof(OutDataType);
+
+            if constexpr(OutputIndex)
+                num_bytes += out_tensor_size * sizeof(IndexDataType);
+
+            float gb_per_sec = num_bytes / 1.E6 / ave_time;
+
+            std::cout << "Perf: " << std::setw(10) << ave_time << " ms, " << gb_per_sec << " GB/s, "
+                      << op_name << std::endl;
+
+            if(ave_time < best_ave_time)
+            {
+                found           = true;
+                best_op_id      = i;
+                best_op_name    = op_name;
+                best_ave_time   = ave_time;
+                best_gb_per_sec = gb_per_sec;
+            }
+        }
+        else
+        {
+            std::cout << op_name << " does not support this problem" << std::endl;
+        }
+    }
+
+    // run the best intance
+    if(found)
+    {
+        std::cout << "Best Perf: " << best_ave_time << " ms, " << best_gb_per_sec << " GB/s, "
+                  << best_op_name << std::endl;
+
+        auto& op_ptr = op_ptrs[best_op_id];
+        std::cout << "Run the best instance without timing: " << op_ptr->GetTypeString()
+                  << std::endl;
+
+        auto argument_ptr = op_ptr->MakeArgumentPointer(
+            static_cast<InDataType*>(in_device_buf.GetDeviceBuffer()),
+            static_cast<OutDataType*>(out_device_buf.GetDeviceBuffer()),
+            static_cast<IndexDataType*>(out_indices_device_buf.GetDeviceBuffer()),
+            in_length,
+            window_spatial_lengths,
+            out_length,
+            in_tensor_stride,
+            out_tensor_stride,
+            out_tensor_stride,
+            window_strides,
+            input_left_pads,
+            input_right_pads,
+            {2, 3});
+
+        auto invoker_ptr = op_ptr->MakeInvokerPointer();
+
+        if(op_ptr->IsSupportedArgument(argument_ptr.get()))
+        {
+            invoker_ptr->Run(argument_ptr.get(), StreamConfig{nullptr, false});
+        }
+
+        std::cout << "Done" << std::endl;
+    }
+
+    return 0;
+}

example/13_pool2d_fwd/pool2d_fwd_common.hpp

@@ -17,115 +17,11 @@
 #include "ck/library/utility/host_tensor.hpp"
 #include "ck/library/utility/host_tensor_generator.hpp"
 #include "ck/library/utility/literals.hpp"
+#include "ck/library/reference_tensor_operation/cpu/reference_pool_fwd.hpp"
 
 template <typename InDataType,
           typename OutDataType,
-          typename AccDataType,
-          typename IndexDataType,
-          ck::ReduceTensorOp ReduceOpId,
-          bool PropagateNan,
-          bool OutputIndex>
-static void pool_host_verify(const Tensor<InDataType>& in,
-                             Tensor<OutDataType>& out,
-                             Tensor<IndexDataType>& out_indices,
-                             const std::array<ck::index_t, 2>& window_spatial_lengths,
-                             const std::array<ck::index_t, 2>& window_strides,
-                             const std::array<ck::index_t, 2>& in_left_pads,
-                             const std::array<ck::index_t, 2>& /*in_right_pads*/)
-{
-    const int32_t reduceLength = window_spatial_lengths[0] * window_spatial_lengths[1];
-
-    using ReduceOperation = typename ck::reduce_binary_operator<ReduceOpId>::opType;
-
-    auto elementwise_ops =
-        ck::reduce_unary_operator<ReduceOpId, true, true>::GetElementwiseOperator(reduceLength);
-
-    auto in_elementwise_op  = std::get<0>(elementwise_ops);
-    auto acc_elementwise_op = std::get<1>(elementwise_ops);
-
-    if constexpr(!OutputIndex)
-    {
-        using Accumulation =
-            ck::detail::AccumulateWithNanCheck<PropagateNan, ReduceOperation, AccDataType>;
-
-        auto f_nchw = [&](auto n, auto c, auto ho, auto wo) {
-            auto accuVal = ReduceOperation::template GetIdentityValue<AccDataType>();
-
-            for(ck::index_t y = 0; y < window_spatial_lengths[0]; ++y)
-            {
-                ck::index_t hi = ho * window_strides[0] + y - in_left_pads[0];
-                for(ck::index_t x = 0; x < window_spatial_lengths[1]; ++x)
-                {
-                    ck::index_t wi = wo * window_strides[1] + x - in_left_pads[1];
-                    if(hi >= 0 && hi < static_cast<ck::index_t>(in.mDesc.GetLengths()[2]) &&
-                       wi >= 0 && wi < static_cast<ck::index_t>(in.mDesc.GetLengths()[3]))
-                    {
-                        AccDataType currVal = static_cast<AccDataType>(in(n, c, hi, wi));
-
-                        in_elementwise_op(currVal, currVal);
-
-                        Accumulation::Calculate(accuVal, currVal);
-                    }
-                }
-            }
-
-            acc_elementwise_op(accuVal, accuVal);
-
-            out(n, c, ho, wo) = accuVal;
-        };
-
-        make_ParallelTensorFunctor(f_nchw,
-                                   out.mDesc.GetLengths()[0],
-                                   out.mDesc.GetLengths()[1],
-                                   out.mDesc.GetLengths()[2],
-                                   out.mDesc.GetLengths()[3])(std::thread::hardware_concurrency());
-    }
-    else
-    {
-        using Accumulation = ck::detail::AccumulateWithIndexAndNanCheck<PropagateNan,
-                                                                        ReduceOperation,
-                                                                        AccDataType,
-                                                                        IndexDataType>;
-        auto f_nchw        = [&](auto n, auto c, auto ho, auto wo) {
-            auto accuVal            = ReduceOperation::template GetIdentityValue<AccDataType>();
-            IndexDataType accuIndex = 0;
-
-            for(ck::index_t y = 0; y < window_spatial_lengths[0]; ++y)
-            {
-                ck::index_t hi = ho * window_strides[0] + y - in_left_pads[0];
-                for(ck::index_t x = 0; x < window_spatial_lengths[1]; ++x)
-                {
-                    ck::index_t wi = wo * window_strides[1] + x - in_left_pads[1];
-                    if(hi >= 0 && hi < in.mDesc.GetLengths()[2] && wi >= 0 &&
-                       wi < in.mDesc.GetLengths()[3])
-                    {
-                        AccDataType currVal     = static_cast<AccDataType>(in(n, c, hi, wi));
-                        IndexDataType currIndex = y * window_spatial_lengths[1] + x;
-
-                        in_elementwise_op(currVal, currVal);
-
-                        Accumulation::Calculate(accuVal, currVal, accuIndex, currIndex);
-                    }
-                }
-            }
-
-            acc_elementwise_op(accuVal, accuVal);
-
-            out(n, c, ho, wo)         = accuVal;
-            out_indices(n, c, ho, wo) = accuIndex;
-        };
-
-        make_ParallelTensorFunctor(f_nchw,
-                                   out.mDesc.GetLengths()[0],
-                                   out.mDesc.GetLengths()[1],
-                                   out.mDesc.GetLengths()[2],
-                                   out.mDesc.GetLengths()[3])(std::thread::hardware_concurrency());
-    };
-}
-
-template <typename InDataType,
-          typename OutDataType,
-          typename AccDataType,
+          typename ComputeDataType,
           typename IndexDataType,
           typename InLayout,
           typename OutLayout,
@@ -150,9 +46,10 @@ bool pool_test(bool do_verification,
 {
     using DevicePoolFwdInstance =
         ck::tensor_operation::device::DevicePool2dFwd_Input_N_Hi_Wi_C_Output_N_Ho_Wo_C<
-            InDataType,  // InDataType
-            OutDataType, // OutDataType
-            AccDataType, // AccDataType
+            InDataType,      // InDataType
+            OutDataType,     // OutDataType
+            IndexDataType,   // IndexDataType
+            ComputeDataType, // ComputeDataType
             ReduceOpId,
             OutputIndex,
             64, // BlockSize
@@ -165,10 +62,10 @@ bool pool_test(bool do_verification,
     const ck::index_t Ho = (Hi + in_left_pad_h + in_right_pad_h - Y) / window_stride_h + 1;
     const ck::index_t Wo = (Wi + in_left_pad_w + in_right_pad_w - X) / window_stride_w + 1;
 
-    const std::array<ck::index_t, 2> window_spatial_lengths{{Y, X}};
-    const std::array<ck::index_t, 2> window_strides{{window_stride_h, window_stride_w}};
-    const std::array<ck::index_t, 2> input_left_pads{{in_left_pad_h, in_left_pad_w}};
-    const std::array<ck::index_t, 2> input_right_pads{{in_right_pad_h, in_right_pad_w}};
+    const std::vector<ck::index_t> window_spatial_lengths{Y, X};
+    const std::vector<ck::index_t> window_strides{window_stride_h, window_stride_w};
+    const std::vector<ck::index_t> input_left_pads{in_left_pad_h, in_left_pad_w};
+    const std::vector<ck::index_t> input_right_pads{in_right_pad_h, in_right_pad_w};
 
     // tensor layout
     auto f_host_tensor_descriptor =
@@ -219,14 +116,16 @@ bool pool_test(bool do_verification,
         static_cast<InDataType*>(in_device_buf.GetDeviceBuffer()),
         static_cast<OutDataType*>(out_device_buf.GetDeviceBuffer()),
         static_cast<IndexDataType*>(out_indices_device_buf.GetDeviceBuffer()),
-        N,
-        C,
-        std::array<ck::index_t, 2>{{Hi, Wi}},
-        std::array<ck::index_t, 2>{{Y, X}},
-        std::array<ck::index_t, 2>{{Ho, Wo}},
+        {N, C, Hi, Wi},
+        {Y, X},
+        {N, C, Ho, Wo},
+        {C * Hi * Wi, 1, Wi * C, C},
+        {C * Ho * Wo, 1, Wo * C, C},
+        {C * Ho * Wo, 1, Wo * C, C},
         window_strides,
         input_left_pads,
-        input_right_pads);
+        input_right_pads,
+        {2, 3});
 
     if(!pool.IsSupportedArgument(argument_ptr.get()))
     {
@@ -252,19 +151,28 @@ bool pool_test(bool do_verification,
 
     if(do_verification)
     {
-        pool_host_verify<InDataType,
-                         OutDataType,
-                         AccDataType,
-                         IndexDataType,
-                         ReduceOpId,
-                         PropagateNan,
-                         OutputIndex>(in_n_c_hi_wi,
-                                      out_n_c_ho_wo_host,
-                                      out_indices_n_c_ho_wo_host,
-                                      window_spatial_lengths,
-                                      window_strides,
-                                      input_left_pads,
-                                      input_right_pads);
+        using ReferencePoolingFwdInstance =
+            ck::tensor_operation::host::ReferencePoolingFwd<4,
+                                                            2,
+                                                            InDataType,
+                                                            OutDataType,
+                                                            ComputeDataType,
+                                                            IndexDataType,
+                                                            ReduceOpId,
+                                                            PropagateNan,
+                                                            OutputIndex>;
+
+        auto ref_pooling          = ReferencePoolingFwdInstance{};
+        auto ref_pooling_invoker  = ref_pooling.MakeInvoker();
+        auto ref_pooling_argument = ref_pooling.MakeArgument(in_n_c_hi_wi,
+                                                             out_n_c_ho_wo_host,
+                                                             out_indices_n_c_ho_wo_host,
+                                                             window_spatial_lengths,
+                                                             window_strides,
+                                                             input_left_pads,
+                                                             input_right_pads);
+
+        ref_pooling_invoker.Run(ref_pooling_argument);
 
         out_device_buf.FromDevice(out_n_c_ho_wo_device.mData.data());
 

example/13_pool2d_fwd/pool2d_fwd_fp16.cpp

@@ -2,7 +2,6 @@
 // Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
 
 #include <iostream>
-#include <cstdlib>
 
 #include "ck/ck.hpp"
 #include "ck/tensor_operation/gpu/device/tensor_layout.hpp"
@@ -10,9 +9,9 @@
 
 #include "pool2d_fwd_common.hpp"
 
-using InDataType  = ck::half_t;
-using OutDataType = ck::half_t;
-using AccDataType = float;
+using InDataType      = ck::half_t;
+using OutDataType     = ck::half_t;
+using ComputeDataType = float;
 
 using IndexDataType = int32_t;
 
@@ -91,7 +90,7 @@ int main(int argc, char* argv[])
 
     bool pass = pool_test<InDataType,
                           OutDataType,
-                          AccDataType,
+                          ComputeDataType,
                           IndexDataType,
                           InLayout,
                           OutLayout,

example/13_pool2d_fwd/pool2d_fwd_fp32.cpp

@@ -2,7 +2,6 @@
 // Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
 
 #include <iostream>
-#include <cstdlib>
 
 #include "ck/ck.hpp"
 #include "ck/utility/reduction_enums.hpp"
@@ -10,9 +9,9 @@
 
 #include "pool2d_fwd_common.hpp"
 
-using InDataType  = float;
-using OutDataType = float;
-using AccDataType = float;
+using InDataType      = float;
+using OutDataType     = float;
+using ComputeDataType = float;
 
 using IndexDataType = int32_t;
 
@@ -91,7 +90,7 @@ int main(int argc, char* argv[])
 
     bool pass = pool_test<InDataType,
                           OutDataType,
-                          AccDataType,
+                          ComputeDataType,
                           IndexDataType,
                           InLayout,
                           OutLayout,

example/48_pool3d_fwd/CMakeLists.txt

+add_example_executable(example_pool3d_fwd_fp16 pool3d_fwd_fp16.cpp)
+

example/48_pool3d_fwd/pool3d_fwd_common.hpp

+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
+#pragma once
+
+#include <iostream>
+
+#include "ck/ck.hpp"
+#include "ck/utility/reduction_enums.hpp"
+#include "ck/utility/reduction_functions_accumulate.hpp"
+#include "ck/tensor_operation/gpu/device/reduction_operator_mapping.hpp"
+#include "ck/tensor_operation/gpu/device/impl/device_pool3d_fwd_ndhwc_ndhwc.hpp"
+#include "ck/tensor_operation/gpu/element/element_wise_operation.hpp"
+
+#include "ck/library/utility/check_err.hpp"
+#include "ck/library/utility/device_memory.hpp"
+#include "ck/library/utility/host_tensor.hpp"
+#include "ck/library/utility/host_tensor_generator.hpp"
+#include "ck/library/utility/literals.hpp"
+#include "ck/library/reference_tensor_operation/cpu/reference_pool_fwd.hpp"
+
+template <typename InDataType,
+          typename OutDataType,
+          typename ComputeDataType,
+          typename IndexDataType,
+          typename InLayout,
+          typename OutLayout,
+          ck::ReduceTensorOp ReduceOpId,
+          bool PropagateNan,
+          bool OutputIndex>
+bool pool3d_test(bool do_verification,
+                 bool time_kernel,
+                 ck::index_t N,
+                 ck::index_t C,
+                 ck::index_t Z,
+                 ck::index_t Y,
+                 ck::index_t X,
+                 ck::index_t Di,
+                 ck::index_t Hi,
+                 ck::index_t Wi,
+                 ck::index_t window_stride_d,
+                 ck::index_t window_stride_h,
+                 ck::index_t window_stride_w,
+                 ck::index_t in_left_pad_d,
+                 ck::index_t in_left_pad_h,
+                 ck::index_t in_left_pad_w,
+                 ck::index_t in_right_pad_d,
+                 ck::index_t in_right_pad_h,
+                 ck::index_t in_right_pad_w)
+{
+    using DevicePoolFwdInstance =
+        ck::tensor_operation::device::DevicePool3dFwd_Input_N_Di_Hi_Wi_C_Output_N_Do_Ho_Wo_C<
+            InDataType,      // InDataType
+            OutDataType,     // OutDataType
+            IndexDataType,   // IndexDataType
+            ComputeDataType, // ComputeDataType
+            ReduceOpId,
+            OutputIndex,
+            64, // BlockSize
+            64, // ReduceMThreadClusterSize
+            1,  // ReduceKThreadClusterSize
+            4,  // ReduceMThreadSliceSize
+            1,  // ReduceKThreadSliceSize
+            4>; // InSrcOutDstVectorSize
+
+    const ck::index_t Do = (Di + in_left_pad_d + in_right_pad_d - Z) / window_stride_d + 1;
+    const ck::index_t Ho = (Hi + in_left_pad_h + in_right_pad_h - Y) / window_stride_h + 1;
+    const ck::index_t Wo = (Wi + in_left_pad_w + in_right_pad_w - X) / window_stride_w + 1;
+
+    const std::vector<ck::index_t> window_spatial_lengths{Z, Y, X};
+    const std::vector<ck::index_t> window_strides{
+        window_stride_d, window_stride_h, window_stride_w};
+    const std::vector<ck::index_t> input_left_pads{in_left_pad_d, in_left_pad_h, in_left_pad_w};
+    const std::vector<ck::index_t> input_right_pads{in_right_pad_d, in_right_pad_h, in_right_pad_w};
+
+    // tensor layout
+    auto f_host_tensor_descriptor = [](std::size_t N_,
+                                       std::size_t C_,
+                                       std::size_t D,
+                                       std::size_t H,
+                                       std::size_t W,
+                                       auto layout) {
+        using namespace ck::literals;
+
+        if constexpr(ck::is_same<decltype(layout), ck::tensor_layout::convolution::NCDHW>::value)
+        {
+            return HostTensorDescriptor({N_, C_, D, H, W},
+                                        {C_ * D * H * W, D * H * W, H * W, W, 1_uz});
+        }
+        else if constexpr(ck::is_same<decltype(layout),
+                                      ck::tensor_layout::convolution::NDHWC>::value)
+        {
+            return HostTensorDescriptor({N_, C_, D, H, W},
+                                        {D * C_ * H * W, 1_uz, C_ * H * W, W * C_, C_});
+        }
+    };
+
+    Tensor<InDataType> in_n_c_di_hi_wi(f_host_tensor_descriptor(N, C, Di, Hi, Wi, InLayout{}));
+    Tensor<OutDataType> out_n_c_do_ho_wo_host(
+        f_host_tensor_descriptor(N, C, Do, Ho, Wo, OutLayout{}));
+    Tensor<IndexDataType> out_indices_n_c_do_ho_wo_host(
+        f_host_tensor_descriptor(N, C, Do, Ho, Wo, OutLayout{}));
+    Tensor<OutDataType> out_n_c_do_ho_wo_device(
+        f_host_tensor_descriptor(N, C, Do, Ho, Wo, OutLayout{}));
+    Tensor<IndexDataType> out_indices_n_c_do_ho_wo_device(
+        f_host_tensor_descriptor(N, C, Do, Ho, Wo, OutLayout{}));
+
+    std::cout << "in_n_c_di_hi_wi: " << in_n_c_di_hi_wi.mDesc << std::endl;
+    std::cout << "out_n_c_do_ho_wo: " << out_n_c_do_ho_wo_host.mDesc << std::endl;
+
+    in_n_c_di_hi_wi.GenerateTensorValue(GeneratorTensor_3<InDataType>{-1.0, 1.0});
+
+    DeviceMem in_device_buf(sizeof(InDataType) * in_n_c_di_hi_wi.mDesc.GetElementSpaceSize());
+    DeviceMem out_device_buf(sizeof(OutDataType) *
+                             out_n_c_do_ho_wo_device.mDesc.GetElementSpaceSize());
+    DeviceMem out_indices_device_buf(sizeof(IndexDataType) *
+                                     out_indices_n_c_do_ho_wo_device.mDesc.GetElementSpaceSize());
+
+    in_device_buf.ToDevice(in_n_c_di_hi_wi.mData.data());
+
+    auto pool         = DevicePoolFwdInstance{};
+    auto invoker_ptr  = pool.MakeInvokerPointer();
+    auto argument_ptr = pool.MakeArgumentPointer(
+        static_cast<InDataType*>(in_device_buf.GetDeviceBuffer()),
+        static_cast<OutDataType*>(out_device_buf.GetDeviceBuffer()),
+        static_cast<IndexDataType*>(out_indices_device_buf.GetDeviceBuffer()),
+        {N, C, Di, Hi, Wi},
+        {Z, Y, X},
+        {N, C, Do, Ho, Wo},
+        {Di * C * Hi * Wi, 1, C * Hi * Wi, Wi * C, C},
+        {Do * C * Ho * Wo, 1, C * Ho * Wo, Wo * C, C},
+        {Do * C * Ho * Wo, 1, C * Ho * Wo, Wo * C, C},
+        window_strides,
+        input_left_pads,
+        input_right_pads,
+        {2, 3, 4});
+
+    if(!pool.IsSupportedArgument(argument_ptr.get()))
+    {
+        throw std::runtime_error("wrong! device_op with the specified compilation parameters does "
+                                 "not support this problem");
+    }
+
+    float ave_time = invoker_ptr->Run(argument_ptr.get(), StreamConfig{nullptr, time_kernel});
+    std::cout << "Perf: " << ave_time << std::endl;
+
+    bool pass = true;
+
+    if(do_verification)
+    {
+        using ReferencePoolingFwdInstance =
+            ck::tensor_operation::host::ReferencePoolingFwd<5,
+                                                            3,
+                                                            InDataType,
+                                                            OutDataType,
+                                                            ComputeDataType,
+                                                            IndexDataType,
+                                                            ReduceOpId,
+                                                            PropagateNan,
+                                                            OutputIndex>;
+
+        auto ref_pooling          = ReferencePoolingFwdInstance{};
+        auto ref_pooling_invoker  = ref_pooling.MakeInvoker();
+        auto ref_pooling_argument = ref_pooling.MakeArgument(in_n_c_di_hi_wi,
+                                                             out_n_c_do_ho_wo_host,
+                                                             out_indices_n_c_do_ho_wo_host,
+                                                             window_spatial_lengths,
+                                                             window_strides,
+                                                             input_left_pads,
+                                                             input_right_pads);
+
+        ref_pooling_invoker.Run(ref_pooling_argument);
+
+        out_device_buf.FromDevice(out_n_c_do_ho_wo_device.mData.data());
+
+        pass = pass && ck::utils::check_err(out_n_c_do_ho_wo_device, out_n_c_do_ho_wo_host);
+
+        if constexpr(OutputIndex)
+        {
+            out_indices_device_buf.FromDevice(out_indices_n_c_do_ho_wo_device.mData.data());
+
+            pass = pass && ck::utils::check_err(out_indices_n_c_do_ho_wo_device,
+                                                out_indices_n_c_do_ho_wo_host);
+        };
+    }
+
+    return (pass);
+};

example/48_pool3d_fwd/pool3d_fwd_fp16.cpp

+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
+
+#include <iostream>
+
+#include "ck/ck.hpp"
+#include "ck/tensor_operation/gpu/device/tensor_layout.hpp"
+#include "ck/utility/reduction_enums.hpp"
+
+#include "pool3d_fwd_common.hpp"
+
+using InDataType      = ck::half_t;
+using OutDataType     = ck::half_t;
+using ComputeDataType = float;
+
+using IndexDataType = int32_t;
+
+using InLayout  = ck::tensor_layout::convolution::NDHWC;
+using OutLayout = ck::tensor_layout::convolution::NDHWC;
+
+#if 1
+static constexpr auto ReduceOpId = ck::ReduceTensorOp::MAX;
+#else
+static constexpr auto ReduceOpId = ck::ReduceTensorOp::AVG;
+#endif
+
+static constexpr bool OutputIndex  = false;
+static constexpr bool PropagateNan = false;
+
+int main()
+{
+    bool do_verification = true;
+    bool time_kernel     = false;
+
+    // Pool shape
+    ck::index_t N               = 2;
+    ck::index_t C               = 32;
+    ck::index_t Z               = 2;
+    ck::index_t Y               = 2;
+    ck::index_t X               = 2;
+    ck::index_t Di              = 30;
+    ck::index_t Hi              = 30;
+    ck::index_t Wi              = 30;
+    ck::index_t window_stride_d = 2;
+    ck::index_t window_stride_h = 2;
+    ck::index_t window_stride_w = 2;
+    ck::index_t in_left_pad_d   = 1;
+    ck::index_t in_left_pad_h   = 1;
+    ck::index_t in_left_pad_w   = 1;
+    ck::index_t in_right_pad_d  = 1;
+    ck::index_t in_right_pad_h  = 1;
+    ck::index_t in_right_pad_w  = 1;
+
+    bool pass = pool3d_test<InDataType,
+                            OutDataType,
+                            ComputeDataType,
+                            IndexDataType,
+                            InLayout,
+                            OutLayout,
+                            ReduceOpId,
+                            PropagateNan,
+                            OutputIndex>(do_verification,
+                                         time_kernel,
+                                         N,
+                                         C,
+                                         Z,
+                                         Y,
+                                         X,
+                                         Di,
+                                         Hi,
+                                         Wi,
+                                         window_stride_d,
+                                         window_stride_h,
+                                         window_stride_w,
+                                         in_left_pad_d,
+                                         in_left_pad_h,
+                                         in_left_pad_w,
+                                         in_right_pad_d,
+                                         in_right_pad_h,
+                                         in_right_pad_w);
+
+    return (pass ? 0 : 1);
+}

include/ck/problem_transform/transform_backward_data_convolution_into_gemm_v4r1_nhwc_kyxc_nhwk.hpp

-// SPDX-License-Identifier: MIT
-// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
-
-#ifndef CK_TRANSFORM_BACKWARD_DATA_CONVOLUTION_INTO_GEMM_V4R1_NHWC_KYXC_NHWK_HPP
-#define CK_TRANSFORM_BACKWARD_DATA_CONVOLUTION_INTO_GEMM_V4R1_NHWC_KYXC_NHWK_HPP
-
-#include "common_header.hpp"
-#include "tensor_descriptor.hpp"
-#include "tensor_descriptor_helper.hpp"
-
-namespace ck {
-
-// Number of GEMMs = YTilde * XTilde
-// GemmM = C
-// GemmN = N * HTildeSlice * WTildeSlice
-// GemmK = K * YDotSlice * XDotSlice
-template <typename... Wei,
-          typename... In,
-          typename... Out,
-          typename ConvStrides,
-          typename ConvDilations,
-          typename InLeftPads,
-          typename InRightPads,
-          index_t IYTildeValue,
-          index_t IXTildeValue,
-          index_t GemmK1Value>
-__host__ __device__ constexpr auto
-transform_backward_data_convolution_into_gemm_v4r1_nhwc_kyxc_nhwk(
-    const TensorDescriptor<Wei...>& wei_k_y_x_c_grid_desc,
-    const TensorDescriptor<Out...>& out_n_ho_wo_k_grid_desc,
-    const TensorDescriptor<In...>& in_n_hi_wi_c_grid_desc,
-    const ConvStrides& conv_strides,
-    const ConvDilations& conv_dilations,
-    const InLeftPads& in_left_pads,
-    const InRightPads& in_right_pads,
-    Number<IYTildeValue>,
-    Number<IXTildeValue>,
-    Number<GemmK1Value>)
-{
-    constexpr auto I0 = Number<0>{};
-    constexpr auto I1 = Number<1>{};
-    constexpr auto I2 = Number<2>{};
-    constexpr auto I3 = Number<3>{};
-
-    constexpr auto GemmK1  = Number<GemmK1Value>{};
-    constexpr auto IYTilde = Number<IYTildeValue>{};
-    constexpr auto IXTilde = Number<IXTildeValue>{};
-
-    const auto N = in_n_hi_wi_c_grid_desc.GetLength(I0);
-    const auto C = in_n_hi_wi_c_grid_desc.GetLength(I3);
-    const auto K = out_n_ho_wo_k_grid_desc.GetLength(I3);
-
-    const auto Hi = in_n_hi_wi_c_grid_desc.GetLength(I1);
-    const auto Wi = in_n_hi_wi_c_grid_desc.GetLength(I2);
-
-    const auto Ho = out_n_ho_wo_k_grid_desc.GetLength(I1);
-    const auto Wo = out_n_ho_wo_k_grid_desc.GetLength(I2);
-
-    const auto Y = wei_k_y_x_c_grid_desc.GetLength(I1);
-    const auto X = wei_k_y_x_c_grid_desc.GetLength(I2);
-
-    const auto ConvStrideH = conv_strides[I0];
-    const auto ConvStrideW = conv_strides[I1];
-
-    const auto ConvDilationH = conv_dilations[I0];
-    const auto ConvDilationW = conv_dilations[I1];
-
-    const auto InLeftPadH = in_left_pads[I0];
-    const auto InLeftPadW = in_left_pads[I1];
-
-    const auto InRightPadH = in_right_pads[I0];
-    const auto InRightPadW = in_right_pads[I1];
-
-    const auto GcdStrideDilationH = math::gcd(ConvStrideH, ConvDilationH);
-    const auto GcdStrideDilationW = math::gcd(ConvStrideW, ConvDilationW);
-
-    const auto YTilde = ConvStrideH / GcdStrideDilationH;
-    const auto XTilde = ConvStrideW / GcdStrideDilationW;
-
-    const auto YDot = math::integer_divide_ceil(Y, YTilde);
-    const auto XDot = math::integer_divide_ceil(X, XTilde);
-
-    const auto HTilde = Ho + math::integer_divide_ceil(ConvDilationH * (Y - I1), ConvStrideH);
-    const auto WTilde = Wo + math::integer_divide_ceil(ConvDilationW * (X - I1), ConvStrideW);
-
-    // only work on HTilde and WTilde that contribute to non-padding area of input tensor
-    const auto IHTildeSliceBegin = math::integer_divide_floor(
-        math::max(I0, InLeftPadH - ConvDilationH * (YTilde - I1)), ConvStrideH);
-    const auto IWTildeSliceBegin = math::integer_divide_floor(
-        math::max(I0, InLeftPadW - ConvDilationW * (XTilde - I1)), ConvStrideW);
-
-    const auto IHTildeSliceEnd =
-        math::min(HTilde, math::integer_divide_ceil(InLeftPadH + Hi - I1, ConvStrideH) + I1);
-    const auto IWTildeSliceEnd =
-        math::min(WTilde, math::integer_divide_ceil(InLeftPadW + Wi - I1, ConvStrideW) + I1);
-
-    const auto HTildeSlice = IHTildeSliceEnd - IHTildeSliceBegin;
-    const auto WTildeSlice = IWTildeSliceEnd - IWTildeSliceBegin;
-
-    // GemmK is different for each GEMM
-    const auto YDotSlice = math::integer_divide_ceil(Y - IYTilde, YTilde);
-    const auto XDotSlice = math::integer_divide_ceil(X - IXTilde, XTilde);
-
-    const auto K1 = GemmK1;
-    const auto K0 = K / K1;
-
-    // weight tensor
-    const auto wei_k_ydot_ytilde_xdot_xtilde_c_grid_desc = transform_tensor_descriptor(
-        wei_k_y_x_c_grid_desc,
-        make_tuple(make_pass_through_transform(K),
-                   make_embed_transform(make_tuple(YDot, YTilde),
-                                        make_tuple(ConvStrideH / GcdStrideDilationH, I1)),
-                   make_embed_transform(make_tuple(XDot, XTilde),
-                                        make_tuple(ConvStrideW / GcdStrideDilationW, I1)),
-                   make_pass_through_transform(C)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}),
-        make_tuple(Sequence<0>{}, Sequence<1, 2>{}, Sequence<3, 4>{}, Sequence<5>{}));
-
-    const auto wei_k0_k1_ydotslice_xdotslice_c_grid_desc =
-        transform_tensor_descriptor(wei_k_ydot_ytilde_xdot_xtilde_c_grid_desc,
-                                    make_tuple(make_unmerge_transform(make_tuple(K0, K1)),
-                                               make_slice_transform(YDot, I0, YDotSlice),
-                                               make_slice_transform(XDot, I0, XDotSlice),
-                                               make_freeze_transform(IYTilde),
-                                               make_freeze_transform(IXTilde),
-                                               make_pass_through_transform(C)),
-                                    make_tuple(Sequence<0>{},
-                                               Sequence<1>{},
-                                               Sequence<3>{},
-                                               Sequence<2>{},
-                                               Sequence<4>{},
-                                               Sequence<5>{}),
-                                    make_tuple(Sequence<0, 1>{},
-                                               Sequence<2>{},
-                                               Sequence<3>{},
-                                               Sequence<>{},
-                                               Sequence<>{},
-                                               Sequence<4>{}));
-
-#if 1
-    const auto wei_gemmk0_gemmm_gemmk1_grid_desc = transform_tensor_descriptor(
-        wei_k0_k1_ydotslice_xdotslice_c_grid_desc,
-        make_tuple(make_merge_transform(make_tuple(YDotSlice, XDotSlice, K0)),
-                   make_pass_through_transform(C),
-                   make_pass_through_transform(K1)),
-        make_tuple(Sequence<2, 3, 0>{}, Sequence<4>{}, Sequence<1>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}));
-#else
-    const auto wei_gemmk0_gemmm_gemmk1_grid_desc = transform_tensor_descriptor(
-        wei_k0_k1_ydotslice_xdotslice_c_grid_desc,
-        make_tuple(make_merge_transform(make_tuple(K0, YDotSlice, XDotSlice)),
-                   make_pass_through_transform(C),
-                   make_pass_through_transform(K1)),
-        make_tuple(Sequence<0, 2, 3>{}, Sequence<4>{}, Sequence<1>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}));
-#endif
-
-    // output tensor
-    // this add padding check
-    const auto out_n_hop_wop_k_grid_desc = transform_tensor_descriptor(
-        out_n_ho_wo_k_grid_desc,
-        make_tuple(make_pass_through_transform(N),
-                   make_pad_transform(Ho, I0, I0),
-                   make_pad_transform(Wo, I0, I0),
-                   make_pass_through_transform(K)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}));
-
-    const auto out_n_ydot_htilde_xdot_wtilde_k_grid_desc = transform_tensor_descriptor(
-        out_n_hop_wop_k_grid_desc,
-        make_tuple(make_pass_through_transform(N),
-                   make_embed_transform(make_tuple(YDot, HTilde),
-                                        make_tuple(-ConvDilationH / GcdStrideDilationH, I1)),
-                   make_embed_transform(make_tuple(XDot, WTilde),
-                                        make_tuple(-ConvDilationW / GcdStrideDilationW, I1)),
-                   make_pass_through_transform(K)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}),
-        make_tuple(Sequence<0>{}, Sequence<1, 2>{}, Sequence<3, 4>{}, Sequence<5>{}));
-
-    const auto out_n_ydotslice_htildeslice_xdotslice_wtildeslice_k0_k1_grid_desc =
-        transform_tensor_descriptor(
-            out_n_ydot_htilde_xdot_wtilde_k_grid_desc,
-            make_tuple(make_pass_through_transform(N),
-                       make_slice_transform(YDot, I0, YDotSlice),
-                       make_slice_transform(HTilde, IHTildeSliceBegin, HTildeSlice),
-                       make_slice_transform(XDot, I0, XDotSlice),
-                       make_slice_transform(WTilde, IWTildeSliceBegin, WTildeSlice),
-                       make_unmerge_transform(make_tuple(K0, K1))),
-            make_tuple(Sequence<0>{},
-                       Sequence<1>{},
-                       Sequence<2>{},
-                       Sequence<3>{},
-                       Sequence<4>{},
-                       Sequence<5>{}),
-            make_tuple(Sequence<0>{},
-                       Sequence<1>{},
-                       Sequence<2>{},
-                       Sequence<3>{},
-                       Sequence<4>{},
-                       Sequence<5, 6>{}));
-
-#if 1
-    const auto out_gemmk0_gemmn_gemmk1_grid_desc = transform_tensor_descriptor(
-        out_n_ydotslice_htildeslice_xdotslice_wtildeslice_k0_k1_grid_desc,
-        make_tuple(make_merge_transform(make_tuple(YDotSlice, XDotSlice, K0)),
-                   make_merge_transform(make_tuple(N, HTildeSlice, WTildeSlice)),
-                   make_pass_through_transform(K1)),
-        make_tuple(Sequence<1, 3, 5>{}, Sequence<0, 2, 4>{}, Sequence<6>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}));
-#else
-    const auto out_gemmk0_gemmn_gemmk1_grid_desc = transform_tensor_descriptor(
-        out_n_ydotslice_htildeslice_xdotslice_wtildeslice_k0_k1_grid_desc,
-        make_tuple(make_merge_transform(make_tuple(K0, YDotSlice, XDotSlice)),
-                   make_merge_transform(make_tuple(N, HTildeSlice, WTildeSlice)),
-                   make_pass_through_transform(K1)),
-        make_tuple(Sequence<5, 1, 3>{}, Sequence<0, 2, 4>{}, Sequence<6>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}));
-#endif
-
-    // input tensor
-    const auto in_n_hip_wip_c_grid_desc = transform_tensor_descriptor(
-        in_n_hi_wi_c_grid_desc,
-        make_tuple(make_pass_through_transform(N),
-                   make_pad_transform(Hi, InLeftPadH, InRightPadH),
-                   make_pad_transform(Wi, InLeftPadW, InRightPadW),
-                   make_pass_through_transform(C)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}));
-
-    const auto in_n_ytilde_htilde_xtilde_wtilde_c_grid_desc = transform_tensor_descriptor(
-        in_n_hip_wip_c_grid_desc,
-        make_tuple(make_pass_through_transform(N),
-                   make_embed_transform(make_tuple(YTilde, HTilde),
-                                        make_tuple(ConvDilationH, ConvStrideH)),
-                   make_embed_transform(make_tuple(XTilde, WTilde),
-                                        make_tuple(ConvDilationW, ConvStrideW)),
-                   make_pass_through_transform(C)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}),
-        make_tuple(Sequence<0>{}, Sequence<1, 2>{}, Sequence<3, 4>{}, Sequence<5>{}));
-
-    const auto in_n_htildeslice_wtildeslice_c_grid_desc = transform_tensor_descriptor(
-        in_n_ytilde_htilde_xtilde_wtilde_c_grid_desc,
-        make_tuple(make_pass_through_transform(N),
-                   make_freeze_transform(IYTilde),
-                   make_slice_transform(HTilde, IHTildeSliceBegin, HTildeSlice),
-                   make_freeze_transform(IXTilde),
-                   make_slice_transform(WTilde, IWTildeSliceBegin, WTildeSlice),
-                   make_pass_through_transform(C)),
-        make_tuple(Sequence<0>{},
-                   Sequence<1>{},
-                   Sequence<2>{},
-                   Sequence<3>{},
-                   Sequence<4>{},
-                   Sequence<5>{}),
-        make_tuple(Sequence<0>{},
-                   Sequence<>{},
-                   Sequence<1>{},
-                   Sequence<>{},
-                   Sequence<2>{},
-                   Sequence<3>{}));
-
-    const auto in_gemmm_gemmn_grid_desc = transform_tensor_descriptor(
-        in_n_htildeslice_wtildeslice_c_grid_desc,
-        make_tuple(make_pass_through_transform(C),
-                   make_merge_transform(make_tuple(N, HTildeSlice, WTildeSlice))),
-        make_tuple(Sequence<3>{}, Sequence<0, 1, 2>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}));
-
-    return make_tuple(wei_gemmk0_gemmm_gemmk1_grid_desc,
-                      out_gemmk0_gemmn_gemmk1_grid_desc,
-                      in_gemmm_gemmn_grid_desc);
-}
-
-} // namespace ck
-#endif

include/ck/problem_transform/transform_backward_data_convolution_into_gemm_v4r1r2_nhwc_kyxc_nhwk.hpp

-// SPDX-License-Identifier: MIT
-// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
-
-#ifndef CK_TRANSFORM_BACKWARD_DATA_CONVOLUTION_INTO_GEMM_V4R1R2_NHWC_KYXC_NHWK_HPP
-#define CK_TRANSFORM_BACKWARD_DATA_CONVOLUTION_INTO_GEMM_V4R1R2_NHWC_KYXC_NHWK_HPP
-
-#include "common_header.hpp"
-#include "tensor_descriptor.hpp"
-#include "tensor_descriptor_helper.hpp"
-
-namespace ck {
-
-// A: out
-// B: wei
-// C: in
-// Number of GEMMs = YTilde * XTilde
-// GemmM = N * HTildeSlice * WTildeSlice
-// GemmN = C
-// GemmK = K * YDotSlice * XDotSlice
-template <typename... Wei,
-          typename... In,
-          typename... Out,
-          typename ConvStrides,
-          typename ConvDilations,
-          typename InLeftPads,
-          typename InRightPads,
-          typename IYTilde,
-          typename IXTilde,
-          index_t GemmK1Value>
-__host__ __device__ constexpr auto
-transform_backward_data_convolution_into_gemm_v4r1r2_nhwc_kyxc_nhwk(
-    const TensorDescriptor<Out...>& out_n_ho_wo_k_grid_desc,
-    const TensorDescriptor<Wei...>& wei_k_y_x_c_grid_desc,
-    const TensorDescriptor<In...>& in_n_hi_wi_c_grid_desc,
-    const ConvStrides& conv_strides,
-    const ConvDilations& conv_dilations,
-    const InLeftPads& in_left_pads,
-    const InRightPads& in_right_pads,
-    IYTilde i_ytilde,
-    IXTilde i_xtilde,
-    Number<GemmK1Value>)
-{
-    constexpr auto I0 = Number<0>{};
-    constexpr auto I1 = Number<1>{};
-    constexpr auto I2 = Number<2>{};
-    constexpr auto I3 = Number<3>{};
-
-    constexpr auto GemmK1 = Number<GemmK1Value>{};
-
-    const auto N = in_n_hi_wi_c_grid_desc.GetLength(I0);
-    const auto C = in_n_hi_wi_c_grid_desc.GetLength(I3);
-    const auto K = out_n_ho_wo_k_grid_desc.GetLength(I3);
-
-    const auto Hi = in_n_hi_wi_c_grid_desc.GetLength(I1);
-    const auto Wi = in_n_hi_wi_c_grid_desc.GetLength(I2);
-
-    const auto Ho = out_n_ho_wo_k_grid_desc.GetLength(I1);
-    const auto Wo = out_n_ho_wo_k_grid_desc.GetLength(I2);
-
-    const auto Y = wei_k_y_x_c_grid_desc.GetLength(I1);
-    const auto X = wei_k_y_x_c_grid_desc.GetLength(I2);
-
-    const auto ConvStrideH = conv_strides[I0];
-    const auto ConvStrideW = conv_strides[I1];
-
-    const auto ConvDilationH = conv_dilations[I0];
-    const auto ConvDilationW = conv_dilations[I1];
-
-    const auto InLeftPadH = in_left_pads[I0];
-    const auto InLeftPadW = in_left_pads[I1];
-
-    const auto InRightPadH = in_right_pads[I0];
-    const auto InRightPadW = in_right_pads[I1];
-
-    const auto GcdStrideDilationH = math::gcd(ConvStrideH, ConvDilationH);
-    const auto GcdStrideDilationW = math::gcd(ConvStrideW, ConvDilationW);
-
-    const auto YTilde = ConvStrideH / GcdStrideDilationH;
-    const auto XTilde = ConvStrideW / GcdStrideDilationW;
-
-    const auto YDot = math::integer_divide_ceil(Y, YTilde);
-    const auto XDot = math::integer_divide_ceil(X, XTilde);
-
-    const auto HTilde = Ho + math::integer_divide_ceil(ConvDilationH * (Y - I1), ConvStrideH);
-    const auto WTilde = Wo + math::integer_divide_ceil(ConvDilationW * (X - I1), ConvStrideW);
-
-    // only work on HTilde and WTilde that contribute to non-padding area of input tensor
-    const auto IHTildeSliceBegin = math::integer_divide_floor(
-        math::max(I0, InLeftPadH - ConvDilationH * (YTilde - I1)), ConvStrideH);
-    const auto IWTildeSliceBegin = math::integer_divide_floor(
-        math::max(I0, InLeftPadW - ConvDilationW * (XTilde - I1)), ConvStrideW);
-
-    const auto IHTildeSliceEnd =
-        math::min(HTilde, math::integer_divide_ceil(InLeftPadH + Hi - I1, ConvStrideH) + I1);
-    const auto IWTildeSliceEnd =
-        math::min(WTilde, math::integer_divide_ceil(InLeftPadW + Wi - I1, ConvStrideW) + I1);
-
-    const auto HTildeSlice = IHTildeSliceEnd - IHTildeSliceBegin;
-    const auto WTildeSlice = IWTildeSliceEnd - IWTildeSliceBegin;
-
-    // GemmK is different for each GEMM
-    const auto YDotSlice = math::integer_divide_ceil(Y - i_ytilde, YTilde);
-    const auto XDotSlice = math::integer_divide_ceil(X - i_xtilde, XTilde);
-
-    const auto K1 = GemmK1;
-    const auto K0 = K / K1;
-
-    // A: output tensor
-    // this add padding check
-    const auto out_n_hop_wop_k_grid_desc = transform_tensor_descriptor(
-        out_n_ho_wo_k_grid_desc,
-        make_tuple(make_pass_through_transform(N),
-                   make_pad_transform(Ho, I0, I0),
-                   make_pad_transform(Wo, I0, I0),
-                   make_pass_through_transform(K)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}));
-
-    const auto out_n_ydot_htilde_xdot_wtilde_k_grid_desc = transform_tensor_descriptor(
-        out_n_hop_wop_k_grid_desc,
-        make_tuple(make_pass_through_transform(N),
-                   make_embed_transform(make_tuple(YDot, HTilde),
-                                        make_tuple(-ConvDilationH / GcdStrideDilationH, I1)),
-                   make_embed_transform(make_tuple(XDot, WTilde),
-                                        make_tuple(-ConvDilationW / GcdStrideDilationW, I1)),
-                   make_pass_through_transform(K)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}),
-        make_tuple(Sequence<0>{}, Sequence<1, 2>{}, Sequence<3, 4>{}, Sequence<5>{}));
-
-    const auto out_n_ydotslice_htildeslice_xdotslice_wtildeslice_k0_k1_grid_desc =
-        transform_tensor_descriptor(
-            out_n_ydot_htilde_xdot_wtilde_k_grid_desc,
-            make_tuple(make_pass_through_transform(N),
-                       make_slice_transform(YDot, I0, YDotSlice),
-                       make_slice_transform(HTilde, IHTildeSliceBegin, HTildeSlice),
-                       make_slice_transform(XDot, I0, XDotSlice),
-                       make_slice_transform(WTilde, IWTildeSliceBegin, WTildeSlice),
-                       make_unmerge_transform(make_tuple(K0, K1))),
-            make_tuple(Sequence<0>{},
-                       Sequence<1>{},
-                       Sequence<2>{},
-                       Sequence<3>{},
-                       Sequence<4>{},
-                       Sequence<5>{}),
-            make_tuple(Sequence<0>{},
-                       Sequence<1>{},
-                       Sequence<2>{},
-                       Sequence<3>{},
-                       Sequence<4>{},
-                       Sequence<5, 6>{}));
-
-#if 1
-    const auto out_gemmk0_gemmm_gemmk1_grid_desc = transform_tensor_descriptor(
-        out_n_ydotslice_htildeslice_xdotslice_wtildeslice_k0_k1_grid_desc,
-        make_tuple(make_merge_transform(make_tuple(YDotSlice, XDotSlice, K0)),
-                   make_merge_transform(make_tuple(N, HTildeSlice, WTildeSlice)),
-                   make_pass_through_transform(K1)),
-        make_tuple(Sequence<1, 3, 5>{}, Sequence<0, 2, 4>{}, Sequence<6>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}));
-#else
-    const auto out_gemmk0_gemmm_gemmk1_grid_desc = transform_tensor_descriptor(
-        out_n_ydotslice_htildeslice_xdotslice_wtildeslice_k0_k1_grid_desc,
-        make_tuple(make_merge_transform(make_tuple(K0, YDotSlice, XDotSlice)),
-                   make_merge_transform(make_tuple(N, HTildeSlice, WTildeSlice)),
-                   make_pass_through_transform(K1)),
-        make_tuple(Sequence<5, 1, 3>{}, Sequence<0, 2, 4>{}, Sequence<6>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}));
-#endif
-
-    // B: weight tensor
-    const auto wei_k_ydot_ytilde_xdot_xtilde_c_grid_desc = transform_tensor_descriptor(
-        wei_k_y_x_c_grid_desc,
-        make_tuple(make_pass_through_transform(K),
-                   make_embed_transform(make_tuple(YDot, YTilde),
-                                        make_tuple(ConvStrideH / GcdStrideDilationH, I1)),
-                   make_embed_transform(make_tuple(XDot, XTilde),
-                                        make_tuple(ConvStrideW / GcdStrideDilationW, I1)),
-                   make_pass_through_transform(C)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}),
-        make_tuple(Sequence<0>{}, Sequence<1, 2>{}, Sequence<3, 4>{}, Sequence<5>{}));
-
-    const auto wei_k0_k1_ydotslice_xdotslice_c_grid_desc =
-        transform_tensor_descriptor(wei_k_ydot_ytilde_xdot_xtilde_c_grid_desc,
-                                    make_tuple(make_unmerge_transform(make_tuple(K0, K1)),
-                                               make_slice_transform(YDot, I0, YDotSlice),
-                                               make_slice_transform(XDot, I0, XDotSlice),
-                                               make_freeze_transform(i_ytilde),
-                                               make_freeze_transform(i_xtilde),
-                                               make_pass_through_transform(C)),
-                                    make_tuple(Sequence<0>{},
-                                               Sequence<1>{},
-                                               Sequence<3>{},
-                                               Sequence<2>{},
-                                               Sequence<4>{},
-                                               Sequence<5>{}),
-                                    make_tuple(Sequence<0, 1>{},
-                                               Sequence<2>{},
-                                               Sequence<3>{},
-                                               Sequence<>{},
-                                               Sequence<>{},
-                                               Sequence<4>{}));
-
-#if 1
-    const auto wei_gemmk0_gemmn_gemmk1_grid_desc = transform_tensor_descriptor(
-        wei_k0_k1_ydotslice_xdotslice_c_grid_desc,
-        make_tuple(make_merge_transform(make_tuple(YDotSlice, XDotSlice, K0)),
-                   make_pass_through_transform(C),
-                   make_pass_through_transform(K1)),
-        make_tuple(Sequence<2, 3, 0>{}, Sequence<4>{}, Sequence<1>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}));
-#else
-    const auto wei_gemmk0_gemmn_gemmk1_grid_desc = transform_tensor_descriptor(
-        wei_k0_k1_ydotslice_xdotslice_c_grid_desc,
-        make_tuple(make_merge_transform(make_tuple(K0, YDotSlice, XDotSlice)),
-                   make_pass_through_transform(C),
-                   make_pass_through_transform(K1)),
-        make_tuple(Sequence<0, 2, 3>{}, Sequence<4>{}, Sequence<1>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}));
-#endif
-
-    // C: input tensor
-    const auto in_n_hip_wip_c_grid_desc = transform_tensor_descriptor(
-        in_n_hi_wi_c_grid_desc,
-        make_tuple(make_pass_through_transform(N),
-                   make_pad_transform(Hi, InLeftPadH, InRightPadH),
-                   make_pad_transform(Wi, InLeftPadW, InRightPadW),
-                   make_pass_through_transform(C)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}));
-
-    const auto in_n_ytilde_htilde_xtilde_wtilde_c_grid_desc = transform_tensor_descriptor(
-        in_n_hip_wip_c_grid_desc,
-        make_tuple(make_pass_through_transform(N),
-                   make_embed_transform(make_tuple(YTilde, HTilde),
-                                        make_tuple(ConvDilationH, ConvStrideH)),
-                   make_embed_transform(make_tuple(XTilde, WTilde),
-                                        make_tuple(ConvDilationW, ConvStrideW)),
-                   make_pass_through_transform(C)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}),
-        make_tuple(Sequence<0>{}, Sequence<1, 2>{}, Sequence<3, 4>{}, Sequence<5>{}));
-
-    const auto in_n_htildeslice_wtildeslice_c_grid_desc = transform_tensor_descriptor(
-        in_n_ytilde_htilde_xtilde_wtilde_c_grid_desc,
-        make_tuple(make_pass_through_transform(N),
-                   make_freeze_transform(i_ytilde),
-                   make_slice_transform(HTilde, IHTildeSliceBegin, HTildeSlice),
-                   make_freeze_transform(i_xtilde),
-                   make_slice_transform(WTilde, IWTildeSliceBegin, WTildeSlice),
-                   make_pass_through_transform(C)),
-        make_tuple(Sequence<0>{},
-                   Sequence<1>{},
-                   Sequence<2>{},
-                   Sequence<3>{},
-                   Sequence<4>{},
-                   Sequence<5>{}),
-        make_tuple(Sequence<0>{},
-                   Sequence<>{},
-                   Sequence<1>{},
-                   Sequence<>{},
-                   Sequence<2>{},
-                   Sequence<3>{}));
-
-    const auto in_gemmm_gemmn_grid_desc = transform_tensor_descriptor(
-        in_n_htildeslice_wtildeslice_c_grid_desc,
-        make_tuple(make_merge_transform(make_tuple(N, HTildeSlice, WTildeSlice)),
-                   make_pass_through_transform(C)),
-        make_tuple(Sequence<0, 1, 2>{}, Sequence<3>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}));
-
-    return make_tuple(out_gemmk0_gemmm_gemmk1_grid_desc,
-                      wei_gemmk0_gemmn_gemmk1_grid_desc,
-                      in_gemmm_gemmn_grid_desc);
-}
-
-// A: out
-// B: wei
-// C: in
-// Number of GEMMs = 1
-// GemmM = N * Ho * Wo
-// GemmN = C
-// GemmK = K
-template <typename... Wei,
-          typename... In,
-          typename... Out,
-          typename ConvStrides,
-          index_t GemmK1Value>
-__host__ __device__ constexpr auto
-transform_backward_data_convolution_into_gemm_v4r1r2_nhwc_kyxc_nhwk_1x1(
-    const TensorDescriptor<Out...>& out_n_ho_wo_k_grid_desc,
-    const TensorDescriptor<Wei...>& /* wei_k_y_x_c_grid_desc */,
-    const TensorDescriptor<In...>& in_n_hi_wi_c_grid_desc,
-    const ConvStrides& conv_strides,
-    Number<GemmK1Value>)
-{
-    constexpr auto I0 = Number<0>{};
-    constexpr auto I1 = Number<1>{};
-    constexpr auto I2 = Number<2>{};
-    constexpr auto I3 = Number<3>{};
-
-    constexpr auto GemmK1 = Number<GemmK1Value>{};
-
-    const auto N = in_n_hi_wi_c_grid_desc.GetLength(I0);
-    const auto C = in_n_hi_wi_c_grid_desc.GetLength(I3);
-    const auto K = out_n_ho_wo_k_grid_desc.GetLength(I3);
-
-    const auto Ho = out_n_ho_wo_k_grid_desc.GetLength(I1);
-    const auto Wo = out_n_ho_wo_k_grid_desc.GetLength(I2);
-
-    const auto ConvStrideH = conv_strides[I0];
-    const auto ConvStrideW = conv_strides[I1];
-
-    const auto K1 = GemmK1;
-    const auto K0 = K / K1;
-
-    // A: output tensor
-    const auto out_gemmk0_gemmm_gemmk1_grid_desc =
-        transform_tensor_descriptor(make_naive_tensor_descriptor_packed(make_tuple(N * Ho * Wo, K)),
-                                    make_tuple(make_pass_through_transform(N * Ho * Wo),
-                                               make_unmerge_transform(make_tuple(K0, K1))),
-                                    make_tuple(Sequence<0>{}, Sequence<1>{}),
-                                    make_tuple(Sequence<1>{}, Sequence<0, 2>{}));
-
-    // B: weight tensor
-    const auto wei_gemmk0_gemmn_gemmk1_grid_desc = transform_tensor_descriptor(
-        make_naive_tensor_descriptor_packed(make_tuple(K, C)),
-        make_tuple(make_unmerge_transform(make_tuple(K0, K1)), make_pass_through_transform(C)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}),
-        make_tuple(Sequence<0, 2>{}, Sequence<1>{}));
-
-    // C: input tensor
-    const auto in_n_y_ho_x_wo_c_grid_desc = transform_tensor_descriptor(
-        in_n_hi_wi_c_grid_desc,
-        make_tuple(make_pass_through_transform(N),
-                   make_embed_transform(make_tuple(I1, Ho), make_tuple(I1, ConvStrideH)),
-                   make_embed_transform(make_tuple(I1, Wo), make_tuple(I1, ConvStrideW)),
-                   make_pass_through_transform(C)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}),
-        make_tuple(Sequence<0>{}, Sequence<1, 2>{}, Sequence<3, 4>{}, Sequence<5>{}));
-
-    const auto in_gemmm_gemmn_grid_desc = transform_tensor_descriptor(
-        in_n_y_ho_x_wo_c_grid_desc,
-        make_tuple(make_freeze_transform(I0),
-                   make_freeze_transform(I0),
-                   make_merge_transform(make_tuple(N, Ho, Wo)),
-                   make_pass_through_transform(C)),
-        make_tuple(Sequence<1>{}, Sequence<3>{}, Sequence<0, 2, 4>{}, Sequence<5>{}),
-        make_tuple(Sequence<>{}, Sequence<>{}, Sequence<0>{}, Sequence<1>{}));
-
-    return make_tuple(out_gemmk0_gemmm_gemmk1_grid_desc,
-                      wei_gemmk0_gemmn_gemmk1_grid_desc,
-                      in_gemmm_gemmn_grid_desc);
-}
-
-} // namespace ck
-#endif

include/ck/problem_transform/transform_backward_weight_convolution_into_gemm_v4r4r2_atomic_nchw_kcyx_nkhw.hpp

-// SPDX-License-Identifier: MIT
-// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
-
-#ifndef CK_TRANSFORM_BACKWARD_WEIGHT_CONVOLUTION_INTO_GEMM_V4R4R2_ATOMIC_NCHW_KCYX_NKHW_HPP
-#define CK_TRANSFORM_BACKWARD_WEIGHT_CONVOLUTION_INTO_GEMM_V4R4R2_ATOMIC_NCHW_KCYX_NKHW_HPP
-
-#include "common_header.hpp"
-#include "tensor_descriptor.hpp"
-#include "tensor_descriptor_helper.hpp"
-
-namespace ck {
-
-// GemmM = K
-// GemmK = N * Ho * Wo
-// GemmN = C * Y * X
-template <typename... Wei,
-          typename... In,
-          typename... Out,
-          typename ConvStrides,
-          typename ConvDilations,
-          typename InLeftPads,
-          typename InRightPads,
-          index_t GemmK1Value,
-          typename GemmKBatchType,
-          typename GemmKPadType>
-__host__ __device__ constexpr auto
-transform_backward_weight_convolution_into_gemm_v4r4r2_atomic_nchw_kcyx_nkhw_pad(
-    const TensorDescriptor<Wei...>& wei_k_c_y_x_grid_desc,
-    const TensorDescriptor<In...>& in_n_c_hi_wi_grid_desc,
-    const TensorDescriptor<Out...>& out_n_k_ho_wo_grid_desc,
-    const ConvStrides& conv_strides,
-    const ConvDilations& conv_dilations,
-    const InLeftPads& in_left_pads,
-    const InRightPads& in_right_pads,
-    Number<GemmK1Value>,
-    GemmKBatchType GemmKBatch,
-    GemmKPadType GemmKPad)
-{
-    constexpr auto I0 = Number<0>{};
-    constexpr auto I1 = Number<1>{};
-    constexpr auto I2 = Number<2>{};
-    constexpr auto I3 = Number<3>{};
-
-    constexpr auto GemmK1 = Number<GemmK1Value>{};
-
-    const auto N = in_n_c_hi_wi_grid_desc.GetLength(I0);
-    const auto C = in_n_c_hi_wi_grid_desc.GetLength(I1);
-    const auto K = out_n_k_ho_wo_grid_desc.GetLength(I1);
-
-    const auto Hi = in_n_c_hi_wi_grid_desc.GetLength(I2);
-    const auto Wi = in_n_c_hi_wi_grid_desc.GetLength(I3);
-
-    const auto Ho = out_n_k_ho_wo_grid_desc.GetLength(I2);
-    const auto Wo = out_n_k_ho_wo_grid_desc.GetLength(I3);
-
-    const auto Y = wei_k_c_y_x_grid_desc.GetLength(I2);
-    const auto X = wei_k_c_y_x_grid_desc.GetLength(I3);
-
-    const auto ConvStrideH = conv_strides[I0];
-    const auto ConvStrideW = conv_strides[I1];
-
-    const auto ConvDilationH = conv_dilations[I0];
-    const auto ConvDilationW = conv_dilations[I1];
-
-    const auto InLeftPadH = in_left_pads[I0];
-    const auto InLeftPadW = in_left_pads[I1];
-
-    const auto InRightPadH = in_right_pads[I0];
-    const auto InRightPadW = in_right_pads[I1];
-
-    const auto GemmM      = K;
-    const auto GemmN      = C * Y * X;
-    const auto GemmKTotal = N * Ho * Wo;
-    const index_t GemmK0  = GemmKPad / (GemmKBatch * GemmK1);
-
-    // A: output tensor
-    const auto out_gemmktotal_gemmm_grid_desc = transform_tensor_descriptor(
-        make_naive_tensor_descriptor_packed(make_tuple(N, K, Ho * Wo)),
-        make_tuple(make_pass_through_transform(K), make_merge_transform(make_tuple(N, Ho * Wo))),
-        make_tuple(Sequence<1>{}, Sequence<0, 2>{}),
-        make_tuple(Sequence<1>{}, Sequence<0>{}));
-
-    const auto out_gemmkpad_gemmm_grid_desc = transform_tensor_descriptor(
-        out_gemmktotal_gemmm_grid_desc,
-        make_tuple(make_right_pad_transform(GemmKTotal, GemmKPad - GemmKTotal),
-                   make_pass_through_transform(GemmM)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}));
-
-    const auto out_gemmkbatch_gemmk0_gemmm_gemmk1_grid_desc = transform_tensor_descriptor(
-        out_gemmkpad_gemmm_grid_desc,
-        make_tuple(make_unmerge_transform(make_tuple(GemmKBatch, GemmK0, GemmK1)),
-                   make_pass_through_transform(GemmM)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}),
-        make_tuple(Sequence<0, 1, 3>{}, Sequence<2>{}));
-
-    // B: input tensor
-    const auto in_n_c_hip_wip_grid_desc = transform_tensor_descriptor(
-        in_n_c_hi_wi_grid_desc,
-        make_tuple(make_pass_through_transform(N),
-                   make_pass_through_transform(C),
-                   make_pad_transform(Hi, InLeftPadH, InRightPadH),
-                   make_pad_transform(Wi, InLeftPadW, InRightPadW)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}));
-
-    const auto in_n_c_y_ho_x_wo_grid_desc = transform_tensor_descriptor(
-        in_n_c_hip_wip_grid_desc,
-        make_tuple(make_pass_through_transform(N),
-                   make_pass_through_transform(C),
-                   make_embed_transform(make_tuple(Y, Ho), make_tuple(ConvDilationH, ConvStrideH)),
-                   make_embed_transform(make_tuple(X, Wo), make_tuple(ConvDilationW, ConvStrideW))),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2, 3>{}, Sequence<4, 5>{}));
-
-    const auto in_gemmktotal_gemmn_grid_desc =
-        transform_tensor_descriptor(in_n_c_y_ho_x_wo_grid_desc,
-                                    make_tuple(make_merge_transform(make_tuple(C, Y, X)),
-                                               make_merge_transform(make_tuple(N, Ho, Wo))),
-                                    make_tuple(Sequence<1, 2, 4>{}, Sequence<0, 3, 5>{}),
-                                    make_tuple(Sequence<1>{}, Sequence<0>{}));
-
-    const auto in_gemmkpad_gemmn_grid_desc = transform_tensor_descriptor(
-        in_gemmktotal_gemmn_grid_desc,
-        make_tuple(make_right_pad_transform(GemmKTotal, GemmKPad - GemmKTotal),
-                   make_pass_through_transform(GemmN)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}));
-
-    const auto in_gemmkbatch_gemmk0_gemmn_gemmk1_grid_desc = transform_tensor_descriptor(
-        in_gemmkpad_gemmn_grid_desc,
-        make_tuple(make_unmerge_transform(make_tuple(GemmKBatch, GemmK0, GemmK1)),
-                   make_pass_through_transform(GemmN)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}),
-        make_tuple(Sequence<0, 1, 3>{}, Sequence<2>{}));
-
-    // C: weight tensor
-    const auto wei_gemmm_gemmn_grid_desc = transform_tensor_descriptor(
-        make_naive_tensor_descriptor_packed(make_tuple(K, C * Y * X)),
-        make_tuple(make_pass_through_transform(K), make_pass_through_transform(C * Y * X)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}));
-
-    return make_tuple(out_gemmkbatch_gemmk0_gemmm_gemmk1_grid_desc,
-                      in_gemmkbatch_gemmk0_gemmn_gemmk1_grid_desc,
-                      wei_gemmm_gemmn_grid_desc);
-}
-
-} // namespace ck
-#endif

include/ck/problem_transform/transform_backward_weight_convolution_into_gemm_v4r4r2_nchw_kcyx_nkhw.hpp

-// SPDX-License-Identifier: MIT
-// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
-
-#ifndef CK_TRANSFORM_BACKWARD_WEIGHT_CONVOLUTION_INTO_GEMM_V4R4R2_NCHW_KCYX_NKHW_HPP
-#define CK_TRANSFORM_BACKWARD_WEIGHT_CONVOLUTION_INTO_GEMM_V4R4R2_NCHW_KCYX_NKHW_HPP
-
-#include "common_header.hpp"
-#include "tensor_descriptor.hpp"
-#include "tensor_descriptor_helper.hpp"
-
-namespace ck {
-
-// GemmM = K
-// GemmK = N * Ho * Wo
-// GemmN = C * Y * X
-template <typename... Wei,
-          typename... In,
-          typename... Out,
-          typename ConvStrides,
-          typename ConvDilations,
-          typename InLeftPads,
-          typename InRightPads,
-          index_t GemmK1Value>
-__host__ __device__ constexpr auto
-transform_backward_weight_convolution_into_gemm_v4r4r2_nchw_kcyx_nkhw_pad(
-    const TensorDescriptor<Wei...>& wei_k_c_y_x_grid_desc,
-    const TensorDescriptor<In...>& in_n_c_hi_wi_grid_desc,
-    const TensorDescriptor<Out...>& out_n_k_ho_wo_grid_desc,
-    const ConvStrides& conv_strides,
-    const ConvDilations& conv_dilations,
-    const InLeftPads& in_left_pads,
-    const InRightPads& in_right_pads,
-    Number<GemmK1Value>)
-{
-    constexpr auto I0 = Number<0>{};
-    constexpr auto I1 = Number<1>{};
-    constexpr auto I2 = Number<2>{};
-    constexpr auto I3 = Number<3>{};
-
-    constexpr auto GemmK1 = Number<GemmK1Value>{};
-
-    const auto N = in_n_c_hi_wi_grid_desc.GetLength(I0);
-    const auto C = in_n_c_hi_wi_grid_desc.GetLength(I1);
-    const auto K = out_n_k_ho_wo_grid_desc.GetLength(I1);
-
-    const auto Hi = in_n_c_hi_wi_grid_desc.GetLength(I2);
-    const auto Wi = in_n_c_hi_wi_grid_desc.GetLength(I3);
-
-    const auto Ho = out_n_k_ho_wo_grid_desc.GetLength(I2);
-    const auto Wo = out_n_k_ho_wo_grid_desc.GetLength(I3);
-
-    const auto Y = wei_k_c_y_x_grid_desc.GetLength(I2);
-    const auto X = wei_k_c_y_x_grid_desc.GetLength(I3);
-
-    const auto ConvStrideH = conv_strides[I0];
-    const auto ConvStrideW = conv_strides[I1];
-
-    const auto ConvDilationH = conv_dilations[I0];
-    const auto ConvDilationW = conv_dilations[I1];
-
-    const auto InLeftPadH = in_left_pads[I0];
-    const auto InLeftPadW = in_left_pads[I1];
-
-    const auto InRightPadH = in_right_pads[I0];
-    const auto InRightPadW = in_right_pads[I1];
-
-    const auto GemmM  = K;
-    const auto GemmN  = C * Y * X;
-    const auto GemmK  = N * Ho * Wo;
-    const auto GemmK0 = GemmK / GemmK1;
-
-    // weight tensor
-    const auto wei_gemmm_gemmn_grid_desc = transform_tensor_descriptor(
-        make_naive_tensor_descriptor_packed(make_tuple(K, C * Y * X)),
-        make_tuple(make_pass_through_transform(K), make_pass_through_transform(C * Y * X)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}));
-
-    // input tensor
-    const auto in_n_c_hip_wip_grid_desc = transform_tensor_descriptor(
-        in_n_c_hi_wi_grid_desc,
-        make_tuple(make_pass_through_transform(N),
-                   make_pass_through_transform(C),
-                   make_pad_transform(Hi, InLeftPadH, InRightPadH),
-                   make_pad_transform(Wi, InLeftPadW, InRightPadW)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}));
-
-    const auto in_n_c_y_ho_x_wo_grid_desc = transform_tensor_descriptor(
-        in_n_c_hip_wip_grid_desc,
-        make_tuple(make_pass_through_transform(N),
-                   make_pass_through_transform(C),
-                   make_embed_transform(make_tuple(Y, Ho), make_tuple(ConvDilationH, ConvStrideH)),
-                   make_embed_transform(make_tuple(X, Wo), make_tuple(ConvDilationW, ConvStrideW))),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2, 3>{}, Sequence<4, 5>{}));
-
-    const auto in_gemmk_gemmn_grid_desc =
-        transform_tensor_descriptor(in_n_c_y_ho_x_wo_grid_desc,
-                                    make_tuple(make_merge_transform(make_tuple(C, Y, X)),
-                                               make_merge_transform(make_tuple(N, Ho, Wo))),
-                                    make_tuple(Sequence<1, 2, 4>{}, Sequence<0, 3, 5>{}),
-                                    make_tuple(Sequence<1>{}, Sequence<0>{}));
-
-    const auto in_gemmk0_gemmn_gemmk1_grid_desc =
-        transform_tensor_descriptor(in_gemmk_gemmn_grid_desc,
-                                    make_tuple(make_unmerge_transform(make_tuple(GemmK0, GemmK1)),
-                                               make_pass_through_transform(GemmN)),
-                                    make_tuple(Sequence<0>{}, Sequence<1>{}),
-                                    make_tuple(Sequence<0, 2>{}, Sequence<1>{}));
-
-    // output tensor
-    const auto out_gemmk_gemmm_grid_desc = transform_tensor_descriptor(
-        make_naive_tensor_descriptor_packed(make_tuple(N, K, Ho * Wo)),
-        make_tuple(make_pass_through_transform(K), make_merge_transform(make_tuple(N, Ho * Wo))),
-        make_tuple(Sequence<1>{}, Sequence<0, 2>{}),
-        make_tuple(Sequence<1>{}, Sequence<0>{}));
-
-    const auto out_gemmk0_gemmm_gemmk1_grid_desc =
-        transform_tensor_descriptor(out_gemmk_gemmm_grid_desc,
-                                    make_tuple(make_unmerge_transform(make_tuple(GemmK0, GemmK1)),
-                                               make_pass_through_transform(GemmM)),
-                                    make_tuple(Sequence<0>{}, Sequence<1>{}),
-                                    make_tuple(Sequence<0, 2>{}, Sequence<1>{}));
-
-    return make_tuple(out_gemmk0_gemmm_gemmk1_grid_desc,
-                      in_gemmk0_gemmn_gemmk1_grid_desc,
-                      wei_gemmm_gemmn_grid_desc);
-}
-
-} // namespace ck
-#endif

include/ck/problem_transform/transform_backward_weight_convolution_into_gemm_v4r4r4_atomic_nhwc_kyxc_nhwk.hpp

-// SPDX-License-Identifier: MIT
-// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
-
-#ifndef CK_TRANSFORM_BACKWARD_WEIGHT_CONVOLUTION_INTO_GEMM_V4R4R4_ATOMIC_NHWC_KYXC_NHWK_HPP
-#define CK_TRANSFORM_BACKWARD_WEIGHT_CONVOLUTION_INTO_GEMM_V4R4R4_ATOMIC_NHWC_KYXC_NHWK_HPP
-
-#include "common_header.hpp"
-#include "tensor_descriptor.hpp"
-#include "tensor_descriptor_helper.hpp"
-
-namespace ck {
-
-// A: in
-// B: wei
-// C: out
-// GemmM = N * Ho * Wo
-// GemmN = K
-// GemmK = Y * X * C
-template <typename... In,
-          typename... Wei,
-          typename... Out,
-          typename ConvStrides,
-          typename ConvDilations,
-          typename InLeftPads,
-          typename InRightPads,
-          index_t GemmK1Value,
-          typename GemmKBatchType,
-          typename GemmKPadType>
-__host__ __device__ constexpr auto
-transform_backward_weight_convolution_into_gemm_v4r4r4_atomic_nhwc_kyxc_nhwk_pad(
-    const TensorDescriptor<In...>& in_n_hi_wi_c_grid_desc,
-    const TensorDescriptor<Wei...>& wei_k_y_x_c_grid_desc,
-    const TensorDescriptor<Out...>& out_n_ho_wo_k_grid_desc,
-    const ConvStrides& conv_strides,
-    const ConvDilations& conv_dilations,
-    const InLeftPads& in_left_pads,
-    const InRightPads& in_right_pads,
-    Number<GemmK1Value>,
-    GemmKBatchType GemmKBatch,
-    GemmKPadType GemmKPad)
-{
-    constexpr auto I0 = Number<0>{};
-    constexpr auto I1 = Number<1>{};
-    constexpr auto I2 = Number<2>{};
-    constexpr auto I3 = Number<3>{};
-
-    constexpr auto GemmK1 = Number<GemmK1Value>{};
-
-    const auto N = in_n_hi_wi_c_grid_desc.GetLength(I0);
-    const auto C = in_n_hi_wi_c_grid_desc.GetLength(I3);
-    const auto K = out_n_ho_wo_k_grid_desc.GetLength(I3);
-
-    const auto Hi = in_n_hi_wi_c_grid_desc.GetLength(I1);
-    const auto Wi = in_n_hi_wi_c_grid_desc.GetLength(I2);
-
-    const auto Ho = out_n_ho_wo_k_grid_desc.GetLength(I1);
-    const auto Wo = out_n_ho_wo_k_grid_desc.GetLength(I2);
-
-    const auto Y = wei_k_y_x_c_grid_desc.GetLength(I1);
-    const auto X = wei_k_y_x_c_grid_desc.GetLength(I2);
-
-    const auto ConvStrideH = conv_strides[I0];
-    const auto ConvStrideW = conv_strides[I1];
-
-    const auto ConvDilationH = conv_dilations[I0];
-    const auto ConvDilationW = conv_dilations[I1];
-
-    const auto InLeftPadH = in_left_pads[I0];
-    const auto InLeftPadW = in_left_pads[I1];
-
-    const auto InRightPadH = in_right_pads[I0];
-    const auto InRightPadW = in_right_pads[I1];
-
-    const auto GemmM      = Y * X * C;
-    const auto GemmN      = K;
-    const auto GemmKTotal = N * Ho * Wo;
-    const index_t GemmK0  = GemmKPad / (GemmKBatch * GemmK1);
-
-    // A: input tensor
-    const auto in_n_hip_wip_c_grid_desc = transform_tensor_descriptor(
-        in_n_hi_wi_c_grid_desc,
-        make_tuple(make_pass_through_transform(N),
-                   make_pad_transform(Hi, InLeftPadH, InRightPadH),
-                   make_pad_transform(Wi, InLeftPadW, InRightPadW),
-                   make_pass_through_transform(C)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}));
-
-    const auto in_n_y_ho_x_wo_c_grid_desc = transform_tensor_descriptor(
-        in_n_hip_wip_c_grid_desc,
-        make_tuple(make_pass_through_transform(N),
-                   make_embed_transform(make_tuple(Y, Ho), make_tuple(ConvDilationH, ConvStrideH)),
-                   make_embed_transform(make_tuple(X, Wo), make_tuple(ConvDilationW, ConvStrideW)),
-                   make_pass_through_transform(C)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}),
-        make_tuple(Sequence<0>{}, Sequence<1, 2>{}, Sequence<3, 4>{}, Sequence<5>{}));
-
-    const auto in_gemmktotal_gemmm_grid_desc =
-        transform_tensor_descriptor(in_n_y_ho_x_wo_c_grid_desc,
-                                    make_tuple(make_merge_transform(make_tuple(Y, X, C)),
-                                               make_merge_transform(make_tuple(N, Ho, Wo))),
-                                    make_tuple(Sequence<1, 3, 5>{}, Sequence<0, 2, 4>{}),
-                                    make_tuple(Sequence<1>{}, Sequence<0>{}));
-
-    const auto in_gemmkpad_gemmm_grid_desc = transform_tensor_descriptor(
-        in_gemmktotal_gemmm_grid_desc,
-        make_tuple(make_right_pad_transform(GemmKTotal, GemmKPad - GemmKTotal),
-                   make_pass_through_transform(GemmM)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}));
-
-    const auto in_gemmkbatch_gemmk0_gemmm_gemmk1_grid_desc = transform_tensor_descriptor(
-        in_gemmkpad_gemmm_grid_desc,
-        make_tuple(make_unmerge_transform(make_tuple(GemmKBatch, GemmK0, GemmK1)),
-                   make_pass_through_transform(GemmM)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}),
-        make_tuple(Sequence<0, 1, 3>{}, Sequence<2>{}));
-
-    // B: output tensor
-    const auto out_gemmktotal_gemmn_grid_desc =
-        make_naive_tensor_descriptor_packed(make_tuple(N * Ho * Wo, K));
-
-    const auto out_gemmkpad_gemmn_grid_desc = transform_tensor_descriptor(
-        out_gemmktotal_gemmn_grid_desc,
-        make_tuple(make_right_pad_transform(GemmKTotal, GemmKPad - GemmKTotal),
-                   make_pass_through_transform(GemmN)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}));
-
-    const auto out_gemmkbatch_gemmk0_gemmn_gemmk1_grid_desc = transform_tensor_descriptor(
-        out_gemmkpad_gemmn_grid_desc,
-        make_tuple(make_unmerge_transform(make_tuple(GemmKBatch, GemmK0, GemmK1)),
-                   make_pass_through_transform(GemmN)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}),
-        make_tuple(Sequence<0, 1, 3>{}, Sequence<2>{}));
-
-    // C: weight tensor
-    const auto wei_gemmm_gemmn_grid_desc = transform_tensor_descriptor(
-        make_naive_tensor_descriptor_packed(make_tuple(K, Y * X * C)),
-        make_tuple(make_pass_through_transform(K), make_pass_through_transform(Y * X * C)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}),
-        make_tuple(Sequence<1>{}, Sequence<0>{}));
-
-    return make_tuple(in_gemmkbatch_gemmk0_gemmm_gemmk1_grid_desc,
-                      out_gemmkbatch_gemmk0_gemmn_gemmk1_grid_desc,
-                      wei_gemmm_gemmn_grid_desc);
-}
-
-} // namespace ck
-#endif

include/ck/problem_transform/transform_backward_weight_convolution_into_gemm_v4r4r4_nhwc_kyxc_nhwk.hpp

-// SPDX-License-Identifier: MIT
-// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
-
-#ifndef CK_TRANSFORM_BACKWARD_WEIGHT_CONVOLUTION_INTO_GEMM_V4R4R4_NHWC_KYXC_NHWK_HPP
-#define CK_TRANSFORM_BACKWARD_WEIGHT_CONVOLUTION_INTO_GEMM_V4R4R4_NHWC_KYXC_NHWK_HPP
-
-#include "common_header.hpp"
-#include "tensor_descriptor.hpp"
-#include "tensor_descriptor_helper.hpp"
-
-namespace ck {
-
-// A: in
-// B: wei
-// C: out
-// GemmM = N * Ho * Wo
-// GemmN = K
-// GemmK = Y * X * C
-template <typename... In,
-          typename... Wei,
-          typename... Out,
-          typename ConvStrides,
-          typename ConvDilations,
-          typename InLeftPads,
-          typename InRightPads,
-          index_t GemmK1Value>
-__host__ __device__ constexpr auto
-transform_backward_weight_convolution_into_gemm_v4r4r4_nhwc_kyxc_nhwk_pad(
-    const TensorDescriptor<In...>& in_n_hi_wi_c_grid_desc,
-    const TensorDescriptor<Wei...>& wei_k_y_x_c_grid_desc,
-    const TensorDescriptor<Out...>& out_n_ho_wo_k_grid_desc,
-    const ConvStrides& conv_strides,
-    const ConvDilations& conv_dilations,
-    const InLeftPads& in_left_pads,
-    const InRightPads& in_right_pads,
-    Number<GemmK1Value>)
-{
-    constexpr auto I0 = Number<0>{};
-    constexpr auto I1 = Number<1>{};
-    constexpr auto I2 = Number<2>{};
-    constexpr auto I3 = Number<3>{};
-
-    constexpr auto GemmK1 = Number<GemmK1Value>{};
-
-    const auto N = in_n_hi_wi_c_grid_desc.GetLength(I0);
-    const auto C = in_n_hi_wi_c_grid_desc.GetLength(I3);
-    const auto K = out_n_ho_wo_k_grid_desc.GetLength(I3);
-
-    const auto Hi = in_n_hi_wi_c_grid_desc.GetLength(I1);
-    const auto Wi = in_n_hi_wi_c_grid_desc.GetLength(I2);
-
-    const auto Ho = out_n_ho_wo_k_grid_desc.GetLength(I1);
-    const auto Wo = out_n_ho_wo_k_grid_desc.GetLength(I2);
-
-    const auto Y = wei_k_y_x_c_grid_desc.GetLength(I1);
-    const auto X = wei_k_y_x_c_grid_desc.GetLength(I2);
-
-    const auto ConvStrideH = conv_strides[I0];
-    const auto ConvStrideW = conv_strides[I1];
-
-    const auto ConvDilationH = conv_dilations[I0];
-    const auto ConvDilationW = conv_dilations[I1];
-
-    const auto InLeftPadH = in_left_pads[I0];
-    const auto InLeftPadW = in_left_pads[I1];
-
-    const auto InRightPadH = in_right_pads[I0];
-    const auto InRightPadW = in_right_pads[I1];
-
-    const auto GemmM  = Y * X * C;
-    const auto GemmN  = K;
-    const auto GemmK  = N * Ho * Wo;
-    const auto GemmK0 = GemmK / GemmK1;
-
-    // A: input tensor
-    const auto in_n_hip_wip_c_grid_desc = transform_tensor_descriptor(
-        in_n_hi_wi_c_grid_desc,
-        make_tuple(make_pass_through_transform(N),
-                   make_pad_transform(Hi, InLeftPadH, InRightPadH),
-                   make_pad_transform(Wi, InLeftPadW, InRightPadW),
-                   make_pass_through_transform(C)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}));
-
-    const auto in_n_y_ho_x_wo_c_grid_desc = transform_tensor_descriptor(
-        in_n_hip_wip_c_grid_desc,
-        make_tuple(make_pass_through_transform(N),
-                   make_embed_transform(make_tuple(Y, Ho), make_tuple(ConvDilationH, ConvStrideH)),
-                   make_embed_transform(make_tuple(X, Wo), make_tuple(ConvDilationW, ConvStrideW)),
-                   make_pass_through_transform(C)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}),
-        make_tuple(Sequence<0>{}, Sequence<1, 2>{}, Sequence<3, 4>{}, Sequence<5>{}));
-
-    const auto in_gemmk_gemmm_grid_desc =
-        transform_tensor_descriptor(in_n_y_ho_x_wo_c_grid_desc,
-                                    make_tuple(make_merge_transform(make_tuple(Y, X, C)),
-                                               make_merge_transform(make_tuple(N, Ho, Wo))),
-                                    make_tuple(Sequence<1, 3, 5>{}, Sequence<0, 2, 4>{}),
-                                    make_tuple(Sequence<1>{}, Sequence<0>{}));
-
-    const auto in_gemmk0_gemmm_gemmk1_grid_desc =
-        transform_tensor_descriptor(in_gemmk_gemmm_grid_desc,
-                                    make_tuple(make_unmerge_transform(make_tuple(GemmK0, GemmK1)),
-                                               make_pass_through_transform(GemmM)),
-                                    make_tuple(Sequence<0>{}, Sequence<1>{}),
-                                    make_tuple(Sequence<0, 2>{}, Sequence<1>{}));
-
-    // B: output tensor
-    const auto out_gemmk_gemmn_grid_desc = transform_tensor_descriptor(
-        make_naive_tensor_descriptor_packed(make_tuple(N * Ho * Wo, K)),
-        make_tuple(make_pass_through_transform(N * Ho * Wo), make_pass_through_transform(K)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}));
-
-    const auto out_gemmk0_gemmn_gemmk1_grid_desc =
-        transform_tensor_descriptor(out_gemmk_gemmn_grid_desc,
-                                    make_tuple(make_unmerge_transform(make_tuple(GemmK0, GemmK1)),
-                                               make_pass_through_transform(GemmN)),
-                                    make_tuple(Sequence<0>{}, Sequence<1>{}),
-                                    make_tuple(Sequence<0, 2>{}, Sequence<1>{}));
-
-    // C: weight tensor
-    const auto wei_gemmm_gemmn_grid_desc = transform_tensor_descriptor(
-        make_naive_tensor_descriptor_packed(make_tuple(K, Y * X * C)),
-        make_tuple(make_pass_through_transform(K), make_pass_through_transform(Y * X * C)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}),
-        make_tuple(Sequence<1>{}, Sequence<0>{}));
-
-    return make_tuple(in_gemmk0_gemmm_gemmk1_grid_desc,
-                      out_gemmk0_gemmn_gemmk1_grid_desc,
-                      wei_gemmm_gemmn_grid_desc);
-}
-
-} // namespace ck
-#endif

include/ck/problem_transform/transform_backward_weight_convolution_into_gemm_v4r4r5_nhwc_kyxc_nhwk.hpp

-// SPDX-License-Identifier: MIT
-// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
-
-#ifndef CK_TRANSFORM_BACKWARD_WEIGHT_CONVOLUTION_INTO_GEMM_V4R4R5_NHWC_KYXC_NHWK_HPP
-#define CK_TRANSFORM_BACKWARD_WEIGHT_CONVOLUTION_INTO_GEMM_V4R4R5_NHWC_KYXC_NHWK_HPP
-
-#include "common_header.hpp"
-#include "tensor_descriptor.hpp"
-#include "tensor_descriptor_helper.hpp"
-
-namespace ck {
-
-// A: out
-// B: in
-// C: wei
-// GemmM = K
-// GemmN = Y * X * C
-// GemmKTotal = N * Ho * Wo
-template <typename... In,
-          typename... Wei,
-          typename... Out,
-          typename ConvStrides,
-          typename ConvDilations,
-          typename InLeftPads,
-          typename InRightPads,
-          index_t GemmK1Value,
-          typename GemmKBatchType,
-          typename GemmKPadType>
-__host__ __device__ constexpr auto
-transform_backward_weight_convolution_into_gemm_v4r4r5_nhwc_kyxc_nhwk_pad(
-    const TensorDescriptor<In...>& in_n_hi_wi_c_grid_desc,
-    const TensorDescriptor<Wei...>& wei_k_y_x_c_grid_desc,
-    const TensorDescriptor<Out...>& out_n_ho_wo_k_grid_desc,
-    const ConvStrides& conv_strides,
-    const ConvDilations& conv_dilations,
-    const InLeftPads& in_left_pads,
-    const InRightPads& in_right_pads,
-    Number<GemmK1Value>,
-    GemmKBatchType GemmKBatch,
-    GemmKPadType GemmKPad)
-{
-    constexpr auto I0 = Number<0>{};
-    constexpr auto I1 = Number<1>{};
-    constexpr auto I2 = Number<2>{};
-    constexpr auto I3 = Number<3>{};
-
-    constexpr auto GemmK1 = Number<GemmK1Value>{};
-
-    const auto N = in_n_hi_wi_c_grid_desc.GetLength(I0);
-    const auto C = in_n_hi_wi_c_grid_desc.GetLength(I3);
-    const auto K = out_n_ho_wo_k_grid_desc.GetLength(I3);
-
-    const auto Hi = in_n_hi_wi_c_grid_desc.GetLength(I1);
-    const auto Wi = in_n_hi_wi_c_grid_desc.GetLength(I2);
-
-    const auto Ho = out_n_ho_wo_k_grid_desc.GetLength(I1);
-    const auto Wo = out_n_ho_wo_k_grid_desc.GetLength(I2);
-
-    const auto Y = wei_k_y_x_c_grid_desc.GetLength(I1);
-    const auto X = wei_k_y_x_c_grid_desc.GetLength(I2);
-
-    const auto ConvStrideH = conv_strides[I0];
-    const auto ConvStrideW = conv_strides[I1];
-
-    const auto ConvDilationH = conv_dilations[I0];
-    const auto ConvDilationW = conv_dilations[I1];
-
-    const auto InLeftPadH = in_left_pads[I0];
-    const auto InLeftPadW = in_left_pads[I1];
-
-    const auto InRightPadH = in_right_pads[I0];
-    const auto InRightPadW = in_right_pads[I1];
-
-    const auto GemmM      = K;
-    const auto GemmN      = Y * X * C;
-    const auto GemmKTotal = N * Ho * Wo;
-    const index_t GemmK0  = GemmKPad / (GemmKBatch * GemmK1);
-
-    // A: output tensor
-    const auto out_gemmktotal_gemmm_grid_desc =
-        make_naive_tensor_descriptor_packed(make_tuple(N * Ho * Wo, K));
-
-    const auto out_gemmkpad_gemmm_grid_desc = transform_tensor_descriptor(
-        out_gemmktotal_gemmm_grid_desc,
-        make_tuple(make_right_pad_transform(GemmKTotal, GemmKPad - GemmKTotal),
-                   make_pass_through_transform(GemmM)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}));
-
-    const auto out_gemmkbatch_gemmk0_gemmm_gemmk1_grid_desc = transform_tensor_descriptor(
-        out_gemmkpad_gemmm_grid_desc,
-        make_tuple(make_unmerge_transform(make_tuple(GemmKBatch, GemmK0, GemmK1)),
-                   make_pass_through_transform(GemmM)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}),
-        make_tuple(Sequence<0, 1, 3>{}, Sequence<2>{}));
-
-    // B: input tensor
-    const auto in_n_hip_wip_c_grid_desc = transform_tensor_descriptor(
-        in_n_hi_wi_c_grid_desc,
-        make_tuple(make_pass_through_transform(N),
-                   make_pad_transform(Hi, InLeftPadH, InRightPadH),
-                   make_pad_transform(Wi, InLeftPadW, InRightPadW),
-                   make_pass_through_transform(C)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}));
-
-    const auto in_n_y_ho_x_wo_c_grid_desc = transform_tensor_descriptor(
-        in_n_hip_wip_c_grid_desc,
-        make_tuple(make_pass_through_transform(N),
-                   make_embed_transform(make_tuple(Y, Ho), make_tuple(ConvDilationH, ConvStrideH)),
-                   make_embed_transform(make_tuple(X, Wo), make_tuple(ConvDilationW, ConvStrideW)),
-                   make_pass_through_transform(C)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}),
-        make_tuple(Sequence<0>{}, Sequence<1, 2>{}, Sequence<3, 4>{}, Sequence<5>{}));
-
-    const auto in_gemmktotal_gemmn_grid_desc =
-        transform_tensor_descriptor(in_n_y_ho_x_wo_c_grid_desc,
-                                    make_tuple(make_merge_transform(make_tuple(Y, X, C)),
-                                               make_merge_transform(make_tuple(N, Ho, Wo))),
-                                    make_tuple(Sequence<1, 3, 5>{}, Sequence<0, 2, 4>{}),
-                                    make_tuple(Sequence<1>{}, Sequence<0>{}));
-
-    const auto in_gemmkpad_gemmn_grid_desc = transform_tensor_descriptor(
-        in_gemmktotal_gemmn_grid_desc,
-        make_tuple(make_right_pad_transform(GemmKTotal, GemmKPad - GemmKTotal),
-                   make_pass_through_transform(GemmN)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}));
-
-    const auto in_gemmkbatch_gemmk0_gemmn_gemmk1_grid_desc = transform_tensor_descriptor(
-        in_gemmkpad_gemmn_grid_desc,
-        make_tuple(make_unmerge_transform(make_tuple(GemmKBatch, GemmK0, GemmK1)),
-                   make_pass_through_transform(GemmN)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}),
-        make_tuple(Sequence<0, 1, 3>{}, Sequence<2>{}));
-
-    // C: weight tensor
-    const auto wei_gemmm_gemmn_grid_desc =
-        make_naive_tensor_descriptor_packed(make_tuple(K, Y * X * C));
-
-    return make_tuple(out_gemmkbatch_gemmk0_gemmm_gemmk1_grid_desc,
-                      in_gemmkbatch_gemmk0_gemmn_gemmk1_grid_desc,
-                      wei_gemmm_gemmn_grid_desc);
-}
-
-} // namespace ck
-#endif

include/ck/problem_transform/transform_forward_convolution_into_gemm_v4r4_nchw_kcyx_nkhw.hpp

-// SPDX-License-Identifier: MIT
-// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
-
-#ifndef CK_TRANSFORM_FORWARD_CONVOLUTION_INTO_GEMM_V4R4_NCHW_KCYX_NKHW_HPP
-#define CK_TRANSFORM_FORWARD_CONVOLUTION_INTO_GEMM_V4R4_NCHW_KCYX_NKHW_HPP
-
-#include "common_header.hpp"
-#include "tensor_descriptor.hpp"
-#include "tensor_descriptor_helper.hpp"
-
-namespace ck {
-
-// GemmM = K
-// GemmN = N * Ho * Wo
-// GemmK = C * Y * X
-template <typename... Wei,
-          typename... In,
-          typename... Out,
-          typename ConvStrides,
-          typename ConvDilations,
-          typename InLeftPads,
-          typename InRightPads>
-__host__ __device__ constexpr auto transform_forward_convolution_into_gemm_v4r4_nchw_kcyx_nkhw_pad(
-    const TensorDescriptor<Wei...>& wei_k_c_y_x_global_desc,
-    const TensorDescriptor<In...>& in_n_c_hi_wi_global_desc,
-    const TensorDescriptor<Out...>& out_n_k_ho_wo_global_desc,
-    const ConvStrides& conv_strides,
-    const ConvDilations& conv_dilations,
-    const InLeftPads& in_left_pads,
-    const InRightPads& in_right_pads)
-{
-    constexpr auto I0 = Number<0>{};
-    constexpr auto I1 = Number<1>{};
-    constexpr auto I2 = Number<2>{};
-    constexpr auto I3 = Number<3>{};
-
-    const auto N = in_n_c_hi_wi_global_desc.GetLength(I0);
-    const auto C = in_n_c_hi_wi_global_desc.GetLength(I1);
-    const auto K = out_n_k_ho_wo_global_desc.GetLength(I1);
-
-    const auto Hi = in_n_c_hi_wi_global_desc.GetLength(I2);
-    const auto Wi = in_n_c_hi_wi_global_desc.GetLength(I3);
-
-    const auto Ho = out_n_k_ho_wo_global_desc.GetLength(I2);
-    const auto Wo = out_n_k_ho_wo_global_desc.GetLength(I3);
-
-    const auto Y = wei_k_c_y_x_global_desc.GetLength(I2);
-    const auto X = wei_k_c_y_x_global_desc.GetLength(I3);
-
-    const auto ConvStrideH = conv_strides[I0];
-    const auto ConvStrideW = conv_strides[I1];
-
-    const auto ConvDilationH = conv_dilations[I0];
-    const auto ConvDilationW = conv_dilations[I1];
-
-    const auto InLeftPadH = in_left_pads[I0];
-    const auto InLeftPadW = in_left_pads[I1];
-
-    const auto InRightPadH = in_right_pads[I0];
-    const auto InRightPadW = in_right_pads[I1];
-
-    // weight tensor
-    const auto wei_gemmk_gemmm_global_desc = transform_tensor_descriptor(
-        make_naive_tensor_descriptor_packed(make_tuple(K, C * Y * X)),
-        make_tuple(make_pass_through_transform(K), make_pass_through_transform(C * Y * X)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}),
-        make_tuple(Sequence<1>{}, Sequence<0>{}));
-
-    // input tensor
-    const auto in_n_c_hip_wip_global_desc = transform_tensor_descriptor(
-        in_n_c_hi_wi_global_desc,
-        make_tuple(make_pass_through_transform(N),
-                   make_pass_through_transform(C),
-                   make_pad_transform(Hi, InLeftPadH, InRightPadH),
-                   make_pad_transform(Wi, InLeftPadW, InRightPadW)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}));
-
-    const auto in_n_c_y_ho_x_wo_global_desc = transform_tensor_descriptor(
-        in_n_c_hip_wip_global_desc,
-        make_tuple(make_pass_through_transform(N),
-                   make_pass_through_transform(C),
-                   make_embed_transform(make_tuple(Y, Ho), make_tuple(ConvDilationH, ConvStrideH)),
-                   make_embed_transform(make_tuple(X, Wo), make_tuple(ConvDilationW, ConvStrideW))),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2, 3>{}, Sequence<4, 5>{}));
-
-    const auto in_gemmk_gemmn_global_desc =
-        transform_tensor_descriptor(in_n_c_y_ho_x_wo_global_desc,
-                                    make_tuple(make_merge_transform(make_tuple(C, Y, X)),
-                                               make_merge_transform(make_tuple(N, Ho, Wo))),
-                                    make_tuple(Sequence<1, 2, 4>{}, Sequence<0, 3, 5>{}),
-                                    make_tuple(Sequence<0>{}, Sequence<1>{}));
-
-    // output tensor
-    const auto out_gemmm_gemmn_global_desc = transform_tensor_descriptor(
-        make_naive_tensor_descriptor_packed(make_tuple(N, K, Ho * Wo)),
-        make_tuple(make_pass_through_transform(K), make_merge_transform(make_tuple(N, Ho * Wo))),
-        make_tuple(Sequence<1>{}, Sequence<0, 2>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}));
-
-    return make_tuple(
-        wei_gemmk_gemmm_global_desc, in_gemmk_gemmn_global_desc, out_gemmm_gemmn_global_desc);
-}
-
-template <typename... Wei,
-          typename... In,
-          typename... Out,
-          typename ConvStrides,
-          typename ConvDilations,
-          typename InLeftPads,
-          typename InRightPads>
-__host__ __device__ constexpr auto
-transform_forward_convolution_into_gemm_v4r4_nchw_kcyx_nkhw_no_pad(
-    const TensorDescriptor<Wei...>& wei_k_c_y_x_global_desc,
-    const TensorDescriptor<In...>& in_n_c_hi_wi_global_desc,
-    const TensorDescriptor<Out...>& out_n_k_ho_wo_global_desc,
-    const ConvStrides& conv_strides,
-    const ConvDilations& conv_dilations,
-    const InLeftPads& in_left_pads,
-    const InRightPads& in_right_pads)
-{
-    constexpr auto I0 = Number<0>{};
-    constexpr auto I1 = Number<1>{};
-    constexpr auto I2 = Number<2>{};
-    constexpr auto I3 = Number<3>{};
-
-    const auto N = in_n_c_hi_wi_global_desc.GetLength(I0);
-    const auto C = in_n_c_hi_wi_global_desc.GetLength(I1);
-    const auto K = out_n_k_ho_wo_global_desc.GetLength(I1);
-
-    const auto Ho = out_n_k_ho_wo_global_desc.GetLength(I2);
-    const auto Wo = out_n_k_ho_wo_global_desc.GetLength(I3);
-
-    const auto Y = wei_k_c_y_x_global_desc.GetLength(I2);
-    const auto X = wei_k_c_y_x_global_desc.GetLength(I3);
-
-    const auto ConvStrideH = conv_strides[I0];
-    const auto ConvStrideW = conv_strides[I1];
-
-    const auto ConvDilationH = conv_dilations[I0];
-    const auto ConvDilationW = conv_dilations[I1];
-
-    const auto InLeftPadH = in_left_pads[I0];
-    const auto InLeftPadW = in_left_pads[I1];
-
-    const auto InRightPadH = in_right_pads[I0];
-    const auto InRightPadW = in_right_pads[I1];
-
-    assert(InLeftPadH == 0 && InLeftPadW == 0 && InRightPadH == 0 && InRightPadW == 0);
-
-    // weight tensor
-    const auto wei_gemmk_gemmm_global_desc = transform_tensor_descriptor(
-        make_naive_tensor_descriptor_packed(make_tuple(K, C * Y * X)),
-        make_tuple(make_pass_through_transform(K), make_pass_through_transform(C * Y * X)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}),
-        make_tuple(Sequence<1>{}, Sequence<0>{}));
-
-    // input tensor
-    const auto in_n_c_y_ho_x_wo_global_desc = transform_tensor_descriptor(
-        in_n_c_hi_wi_global_desc,
-        make_tuple(make_pass_through_transform(N),
-                   make_pass_through_transform(C),
-                   make_embed_transform(make_tuple(Y, Ho), make_tuple(ConvDilationH, ConvStrideH)),
-                   make_embed_transform(make_tuple(X, Wo), make_tuple(ConvDilationW, ConvStrideW))),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2, 3>{}, Sequence<4, 5>{}));
-
-    const auto in_gemmk_gemmn_global_desc =
-        transform_tensor_descriptor(in_n_c_y_ho_x_wo_global_desc,
-                                    make_tuple(make_merge_transform(make_tuple(C, Y, X)),
-                                               make_merge_transform(make_tuple(N, Ho, Wo))),
-                                    make_tuple(Sequence<1, 2, 4>{}, Sequence<0, 3, 5>{}),
-                                    make_tuple(Sequence<0>{}, Sequence<1>{}));
-
-    // output tensor
-    const auto out_gemmm_gemmn_global_desc = transform_tensor_descriptor(
-        make_naive_tensor_descriptor_packed(make_tuple(N, K, Ho * Wo)),
-        make_tuple(make_pass_through_transform(K), make_merge_transform(make_tuple(N, Ho * Wo))),
-        make_tuple(Sequence<1>{}, Sequence<0, 2>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}));
-
-    return make_tuple(
-        wei_gemmk_gemmm_global_desc, in_gemmk_gemmn_global_desc, out_gemmm_gemmn_global_desc);
-}
-
-template <typename... Wei,
-          typename... In,
-          typename... Out,
-          typename ConvStrides,
-          typename ConvDilations,
-          typename InLeftPads,
-          typename InRightPads>
-__host__ __device__ constexpr auto transform_forward_convolution_into_gemm_v4r4_nchw_kcyx_nkhw_1x1(
-    const TensorDescriptor<Wei...>& wei_k_c_y_x_global_desc,
-    const TensorDescriptor<In...>& in_n_c_hi_wi_global_desc,
-    const TensorDescriptor<Out...>& out_n_k_ho_wo_global_desc,
-    const ConvStrides& conv_strides,
-    const ConvDilations& conv_dilations,
-    const InLeftPads& in_left_pads,
-    const InRightPads& in_right_pads)
-{
-    constexpr auto I0 = Number<0>{};
-    constexpr auto I1 = Number<1>{};
-    constexpr auto I2 = Number<2>{};
-    constexpr auto I3 = Number<3>{};
-
-    const auto N = in_n_c_hi_wi_global_desc.GetLength(I0);
-    const auto C = in_n_c_hi_wi_global_desc.GetLength(I1);
-    const auto K = out_n_k_ho_wo_global_desc.GetLength(I1);
-
-    const auto Ho = out_n_k_ho_wo_global_desc.GetLength(I2);
-    const auto Wo = out_n_k_ho_wo_global_desc.GetLength(I3);
-
-    const auto Y = wei_k_c_y_x_global_desc.GetLength(I2);
-    const auto X = wei_k_c_y_x_global_desc.GetLength(I3);
-
-    const auto ConvStrideH = conv_strides[I0];
-    const auto ConvStrideW = conv_strides[I1];
-
-    const auto ConvDilationH = conv_dilations[I0];
-    const auto ConvDilationW = conv_dilations[I1];
-
-    const auto InLeftPadH = in_left_pads[I0];
-    const auto InLeftPadW = in_left_pads[I1];
-
-    const auto InRightPadH = in_right_pads[I0];
-    const auto InRightPadW = in_right_pads[I1];
-
-    assert(Y == 1 && X == 1 && ConvStrideH == 1 && ConvStrideW == 1 && ConvDilationH == 1 &&
-           ConvDilationW == 1 && InLeftPadH == 0 && InLeftPadW == 0 && InRightPadH == 0 &&
-           InRightPadW == 0);
-
-    // weight tensor
-    const auto wei_gemmk_gemmm_global_desc = transform_tensor_descriptor(
-        make_naive_tensor_descriptor_packed(make_tuple(K, C)),
-        make_tuple(make_pass_through_transform(K), make_pass_through_transform(C)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}),
-        make_tuple(Sequence<1>{}, Sequence<0>{}));
-
-    // input tensor
-    const auto in_gemmk_gemmn_global_desc = transform_tensor_descriptor(
-        in_n_c_hi_wi_global_desc,
-        make_tuple(make_pass_through_transform(C), make_merge_transform(make_tuple(N, Ho, Wo))),
-        make_tuple(Sequence<1>{}, Sequence<0, 2, 3>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}));
-
-    // output tensor
-    const auto out_gemmm_gemmn_global_desc = transform_tensor_descriptor(
-        make_naive_tensor_descriptor_packed(make_tuple(N, K, Ho * Wo)),
-        make_tuple(make_pass_through_transform(K), make_merge_transform(make_tuple(N, Ho * Wo))),
-        make_tuple(Sequence<1>{}, Sequence<0, 2>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}));
-
-    return make_tuple(
-        wei_gemmk_gemmm_global_desc, in_gemmk_gemmn_global_desc, out_gemmm_gemmn_global_desc);
-}
-
-} // namespace ck
-#endif

include/ck/problem_transform/transform_forward_convolution_into_gemm_v4r4_nhwc_kyxc_nhwk.hpp

-// SPDX-License-Identifier: MIT
-// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
-
-#ifndef CK_TRANSFORM_FORWARD_CONVOLUTION_INTO_GEMM_V4R4_NHWC_KYXC_NHWK_HPP
-#define CK_TRANSFORM_FORWARD_CONVOLUTION_INTO_GEMM_V4R4_NHWC_KYXC_NHWK_HPP
-
-#include "common_header.hpp"
-#include "tensor_descriptor.hpp"
-#include "tensor_descriptor_helper.hpp"
-
-namespace ck {
-
-// GemmM = K
-// GemmN = N * Ho * Wo
-// GemmK = C * Y * X
-template <typename... Wei,
-          typename... In,
-          typename... Out,
-          typename ConvStrides,
-          typename ConvDilations,
-          typename InLeftPads,
-          typename InRightPads>
-__host__ __device__ constexpr auto transform_forward_convolution_into_gemm_v4r4_nhwc_kyxc_nhwk_pad(
-    const TensorDescriptor<Wei...>& wei_k_y_x_c_grid_desc,
-    const TensorDescriptor<In...>& in_n_hi_wi_c_grid_desc,
-    const TensorDescriptor<Out...>& out_n_ho_wo_k_grid_desc,
-    const ConvStrides& conv_strides,
-    const ConvDilations& conv_dilations,
-    const InLeftPads& in_left_pads,
-    const InRightPads& in_right_pads)
-{
-    constexpr auto I0 = Number<0>{};
-    constexpr auto I1 = Number<1>{};
-    constexpr auto I2 = Number<2>{};
-    constexpr auto I3 = Number<3>{};
-
-    const auto N = in_n_hi_wi_c_grid_desc.GetLength(I0);
-    const auto C = in_n_hi_wi_c_grid_desc.GetLength(I3);
-    const auto K = out_n_ho_wo_k_grid_desc.GetLength(I3);
-
-    const auto Hi = in_n_hi_wi_c_grid_desc.GetLength(I1);
-    const auto Wi = in_n_hi_wi_c_grid_desc.GetLength(I2);
-
-    const auto Ho = out_n_ho_wo_k_grid_desc.GetLength(I1);
-    const auto Wo = out_n_ho_wo_k_grid_desc.GetLength(I2);
-
-    const auto Y = wei_k_y_x_c_grid_desc.GetLength(I1);
-    const auto X = wei_k_y_x_c_grid_desc.GetLength(I2);
-
-    const auto ConvStrideH = conv_strides[I0];
-    const auto ConvStrideW = conv_strides[I1];
-
-    const auto ConvDilationH = conv_dilations[I0];
-    const auto ConvDilationW = conv_dilations[I1];
-
-    const auto InLeftPadH = in_left_pads[I0];
-    const auto InLeftPadW = in_left_pads[I1];
-
-    const auto InRightPadH = in_right_pads[I0];
-    const auto InRightPadW = in_right_pads[I1];
-
-    // weight tensor
-    const auto wei_gemmk_gemmm_grid_desc = transform_tensor_descriptor(
-        make_naive_tensor_descriptor_packed(make_tuple(K, Y * X * C)),
-        make_tuple(make_pass_through_transform(K), make_pass_through_transform(Y * X * C)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}),
-        make_tuple(Sequence<1>{}, Sequence<0>{}));
-
-    // input tensor
-    const auto in_n_hip_wip_c_grid_desc = transform_tensor_descriptor(
-        in_n_hi_wi_c_grid_desc,
-        make_tuple(make_pass_through_transform(N),
-                   make_pad_transform(Hi, InLeftPadH, InRightPadH),
-                   make_pad_transform(Wi, InLeftPadW, InRightPadW),
-                   make_pass_through_transform(C)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}));
-
-    const auto in_n_y_ho_x_wo_c_grid_desc = transform_tensor_descriptor(
-        in_n_hip_wip_c_grid_desc,
-        make_tuple(make_pass_through_transform(N),
-                   make_embed_transform(make_tuple(Y, Ho), make_tuple(ConvDilationH, ConvStrideH)),
-                   make_embed_transform(make_tuple(X, Wo), make_tuple(ConvDilationW, ConvStrideW)),
-                   make_pass_through_transform(C)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}),
-        make_tuple(Sequence<0>{}, Sequence<1, 2>{}, Sequence<3, 4>{}, Sequence<5>{}));
-
-    const auto in_gemmk_gemmn_grid_desc =
-        transform_tensor_descriptor(in_n_y_ho_x_wo_c_grid_desc,
-                                    make_tuple(make_merge_transform(make_tuple(Y, X, C)),
-                                               make_merge_transform(make_tuple(N, Ho, Wo))),
-                                    make_tuple(Sequence<1, 3, 5>{}, Sequence<0, 2, 4>{}),
-                                    make_tuple(Sequence<0>{}, Sequence<1>{}));
-
-    // output tensor
-    const auto out_gemmm_gemmn_grid_desc = transform_tensor_descriptor(
-        make_naive_tensor_descriptor_packed(make_tuple(N * Ho * Wo, K)),
-        make_tuple(make_pass_through_transform(N * Ho * Wo), make_pass_through_transform(K)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}),
-        make_tuple(Sequence<1>{}, Sequence<0>{}));
-
-    return make_tuple(
-        wei_gemmk_gemmm_grid_desc, in_gemmk_gemmn_grid_desc, out_gemmm_gemmn_grid_desc);
-}
-
-template <typename... Wei,
-          typename... In,
-          typename... Out,
-          typename ConvStrides,
-          typename ConvDilations,
-          typename InLeftPads,
-          typename InRightPads>
-__host__ __device__ constexpr auto transform_forward_convolution_into_gemm_v4r4_nhwc_kyxc_nhwk_1x1(
-    const TensorDescriptor<Wei...>& wei_k_y_x_c_grid_desc,
-    const TensorDescriptor<In...>& in_n_hi_wi_c_grid_desc,
-    const TensorDescriptor<Out...>& out_n_ho_wo_k_grid_desc,
-    const ConvStrides& conv_strides,
-    const ConvDilations& conv_dilations,
-    const InLeftPads& in_left_pads,
-    const InRightPads& in_right_pads)
-{
-    constexpr auto I0 = Number<0>{};
-    constexpr auto I1 = Number<1>{};
-    constexpr auto I2 = Number<2>{};
-    constexpr auto I3 = Number<3>{};
-
-    const auto N = in_n_hi_wi_c_grid_desc.GetLength(I0);
-    const auto C = in_n_hi_wi_c_grid_desc.GetLength(I3);
-    const auto K = out_n_ho_wo_k_grid_desc.GetLength(I3);
-
-    const auto Ho = out_n_ho_wo_k_grid_desc.GetLength(I1);
-    const auto Wo = out_n_ho_wo_k_grid_desc.GetLength(I2);
-
-    const auto Y = wei_k_y_x_c_grid_desc.GetLength(I1);
-    const auto X = wei_k_y_x_c_grid_desc.GetLength(I2);
-
-    const auto ConvStrideH = conv_strides[I0];
-    const auto ConvStrideW = conv_strides[I1];
-
-    const auto ConvDilationH = conv_dilations[I0];
-    const auto ConvDilationW = conv_dilations[I1];
-
-    const auto InLeftPadH = in_left_pads[I0];
-    const auto InLeftPadW = in_left_pads[I1];
-
-    const auto InRightPadH = in_right_pads[I0];
-    const auto InRightPadW = in_right_pads[I1];
-
-    assert(Y == 1 && X == 1 && ConvStrideH == 1 && ConvStrideW == 1 && ConvDilationH == 1 &&
-           ConvDilationW == 1 && InLeftPadH == 0 && InLeftPadW == 0 && InRightPadH == 0 &&
-           InRightPadW == 0);
-
-    // weight tensor
-    const auto wei_gemmk_gemmm_grid_desc = transform_tensor_descriptor(
-        make_naive_tensor_descriptor_packed(make_tuple(K, C)),
-        make_tuple(make_pass_through_transform(K), make_pass_through_transform(C)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}),
-        make_tuple(Sequence<1>{}, Sequence<0>{}));
-
-    // input tensor
-    const auto in_gemmk_gemmn_grid_desc = transform_tensor_descriptor(
-        make_naive_tensor_descriptor_packed(make_tuple(N * Ho * Wo, C)),
-        make_tuple(make_pass_through_transform(N * Ho * Wo), make_pass_through_transform(C)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}),
-        make_tuple(Sequence<1>{}, Sequence<0>{}));
-
-    // output tensor
-    const auto out_gemmm_gemmn_grid_desc = transform_tensor_descriptor(
-        make_naive_tensor_descriptor_packed(make_tuple(N * Ho * Wo, K)),
-        make_tuple(make_pass_through_transform(N * Ho * Wo), make_pass_through_transform(K)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}),
-        make_tuple(Sequence<1>{}, Sequence<0>{}));
-
-    return make_tuple(
-        wei_gemmk_gemmm_grid_desc, in_gemmk_gemmn_grid_desc, out_gemmm_gemmn_grid_desc);
-}
-
-} // namespace ck
-#endif

include/ck/problem_transform/transform_forward_convolution_into_gemm_v4r4r2_nchw_kcyx_nkhw.hpp

-// SPDX-License-Identifier: MIT
-// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
-
-#ifndef CK_TRANSFORM_FORWARD_CONVOLUTION_INTO_GEMM_V4R4R2_NCHW_KCYX_NKHW_HPP
-#define CK_TRANSFORM_FORWARD_CONVOLUTION_INTO_GEMM_V4R4R2_NCHW_KCYX_NKHW_HPP
-
-#include "common_header.hpp"
-#include "tensor_descriptor.hpp"
-#include "tensor_descriptor_helper.hpp"
-
-namespace ck {
-
-// GemmM = K
-// GemmN = N * Ho * Wo
-// GemmK = C * Y * X
-template <typename... Wei,
-          typename... In,
-          typename... Out,
-          typename ConvStrides,
-          typename ConvDilations,
-          typename InLeftPads,
-          typename InRightPads,
-          index_t GemmK1Value>
-__host__ __device__ constexpr auto
-transform_forward_convolution_into_gemm_v4r4r2_nchw_kcyx_nkhw_pad(
-    const TensorDescriptor<Wei...>& wei_k_c_y_x_grid_desc,
-    const TensorDescriptor<In...>& in_n_c_hi_wi_grid_desc,
-    const TensorDescriptor<Out...>& out_n_k_ho_wo_grid_desc,
-    const ConvStrides& conv_strides,
-    const ConvDilations& conv_dilations,
-    const InLeftPads& in_left_pads,
-    const InRightPads& in_right_pads,
-    Number<GemmK1Value>)
-{
-    constexpr auto I0 = Number<0>{};
-    constexpr auto I1 = Number<1>{};
-    constexpr auto I2 = Number<2>{};
-    constexpr auto I3 = Number<3>{};
-
-    constexpr auto GemmK1 = Number<GemmK1Value>{};
-
-    const auto N = in_n_c_hi_wi_grid_desc.GetLength(I0);
-    const auto C = in_n_c_hi_wi_grid_desc.GetLength(I1);
-    const auto K = out_n_k_ho_wo_grid_desc.GetLength(I1);
-
-    const auto Hi = in_n_c_hi_wi_grid_desc.GetLength(I2);
-    const auto Wi = in_n_c_hi_wi_grid_desc.GetLength(I3);
-
-    const auto Ho = out_n_k_ho_wo_grid_desc.GetLength(I2);
-    const auto Wo = out_n_k_ho_wo_grid_desc.GetLength(I3);
-
-    const auto Y = wei_k_c_y_x_grid_desc.GetLength(I2);
-    const auto X = wei_k_c_y_x_grid_desc.GetLength(I3);
-
-    const auto ConvStrideH = conv_strides[I0];
-    const auto ConvStrideW = conv_strides[I1];
-
-    const auto ConvDilationH = conv_dilations[I0];
-    const auto ConvDilationW = conv_dilations[I1];
-
-    const auto InLeftPadH = in_left_pads[I0];
-    const auto InLeftPadW = in_left_pads[I1];
-
-    const auto InRightPadH = in_right_pads[I0];
-    const auto InRightPadW = in_right_pads[I1];
-
-    const auto GemmM  = K;
-    const auto GemmN  = N * Ho * Wo;
-    const auto GemmK  = C * Y * X;
-    const auto GemmK0 = GemmK / GemmK1;
-
-    // weight tensor
-    const auto wei_gemmk_gemmm_grid_desc = transform_tensor_descriptor(
-        make_naive_tensor_descriptor_packed(make_tuple(K, C * Y * X)),
-        make_tuple(make_pass_through_transform(K), make_pass_through_transform(C * Y * X)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}),
-        make_tuple(Sequence<1>{}, Sequence<0>{}));
-
-    const auto wei_gemmk0_gemmm_gemmk1_grid_desc =
-        transform_tensor_descriptor(wei_gemmk_gemmm_grid_desc,
-                                    make_tuple(make_unmerge_transform(make_tuple(GemmK0, GemmK1)),
-                                               make_pass_through_transform(GemmM)),
-                                    make_tuple(Sequence<0>{}, Sequence<1>{}),
-                                    make_tuple(Sequence<0, 2>{}, Sequence<1>{}));
-
-    // input tensor
-    const auto in_n_c_hip_wip_grid_desc = transform_tensor_descriptor(
-        in_n_c_hi_wi_grid_desc,
-        make_tuple(make_pass_through_transform(N),
-                   make_pass_through_transform(C),
-                   make_pad_transform(Hi, InLeftPadH, InRightPadH),
-                   make_pad_transform(Wi, InLeftPadW, InRightPadW)),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}));
-
-    const auto in_n_c_y_ho_x_wo_grid_desc = transform_tensor_descriptor(
-        in_n_c_hip_wip_grid_desc,
-        make_tuple(make_pass_through_transform(N),
-                   make_pass_through_transform(C),
-                   make_embed_transform(make_tuple(Y, Ho), make_tuple(ConvDilationH, ConvStrideH)),
-                   make_embed_transform(make_tuple(X, Wo), make_tuple(ConvDilationW, ConvStrideW))),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2, 3>{}, Sequence<4, 5>{}));
-
-    const auto in_gemmk_gemmn_grid_desc =
-        transform_tensor_descriptor(in_n_c_y_ho_x_wo_grid_desc,
-                                    make_tuple(make_merge_transform(make_tuple(C, Y, X)),
-                                               make_merge_transform(make_tuple(N, Ho, Wo))),
-                                    make_tuple(Sequence<1, 2, 4>{}, Sequence<0, 3, 5>{}),
-                                    make_tuple(Sequence<0>{}, Sequence<1>{}));
-
-    const auto in_gemmk0_gemmn_gemmk1_grid_desc =
-        transform_tensor_descriptor(in_gemmk_gemmn_grid_desc,
-                                    make_tuple(make_unmerge_transform(make_tuple(GemmK0, GemmK1)),
-                                               make_pass_through_transform(GemmN)),
-                                    make_tuple(Sequence<0>{}, Sequence<1>{}),
-                                    make_tuple(Sequence<0, 2>{}, Sequence<1>{}));
-
-    // output tensor
-    const auto out_gemmm_gemmn_grid_desc = transform_tensor_descriptor(
-        make_naive_tensor_descriptor_packed(make_tuple(N, K, Ho * Wo)),
-        make_tuple(make_pass_through_transform(K), make_merge_transform(make_tuple(N, Ho * Wo))),
-        make_tuple(Sequence<1>{}, Sequence<0, 2>{}),
-        make_tuple(Sequence<0>{}, Sequence<1>{}));
-
-    return make_tuple(wei_gemmk0_gemmm_gemmk1_grid_desc,
-                      in_gemmk0_gemmn_gemmk1_grid_desc,
-                      out_gemmm_gemmn_grid_desc);
-}
-
-} // namespace ck
-#endif