bring back col2im test

composable_kernel/include/kernel_algorithm/gridwise_col2im_eb_nchw.hpp

+#ifndef CK_GRIDWISE_COL2IM_EB_NCHW_HPP
+#define CK_GRIDWISE_COL2IM_EB_NCHW_HPP
+
+#include "common_header.hpp"
+#include "tensor_descriptor.hpp"
+#include "tensor_descriptor_helper.hpp"
+#include "blockwise_generic_tensor_slice_copy.hpp"
+
+namespace ck {
+
+// B = merge(N, Ho, Wo)
+template <index_t GridSize,
+          index_t BlockSize,
+          typename Float,
+          typename ColGlobalDesc,
+          typename ImgGlobalDesc,
+          typename FilterSizes,
+          typename OutputSizes,
+          typename ConvStrides,
+          typename ConvDilations,
+          typename LeftPads,
+          typename RightPads,
+          index_t EPerBlock,
+          index_t BPerBlock,
+          typename BlockCopySubLengths_E_B,
+          typename BlockCopyClusterLengths_E_B,
+          typename BlockCopyThreadClusterArrangeOrder,
+          typename BlockCopySrcAccessOrder,
+          typename BlockCopyDstAccessOrder,
+          index_t BlockCopyDataPerAccess_B>
+struct GridwiseCol2Im_eb_nchw
+{
+    __device__ void Run(const Float* const __restrict__ p_col_global,
+                        Float* const __restrict__ p_img_global) const
+    {
+        constexpr auto col_e_b_global_desc       = ColGlobalDesc{};
+        constexpr auto img_n_c_hi_wi_global_desc = ImgGlobalDesc{};
+
+        constexpr index_t N  = img_n_c_hi_wi_global_desc.GetLengths()[0];
+        constexpr index_t C  = img_n_c_hi_wi_global_desc.GetLengths()[1];
+        constexpr index_t Hi = img_n_c_hi_wi_global_desc.GetLengths()[2];
+        constexpr index_t Wi = img_n_c_hi_wi_global_desc.GetLengths()[3];
+
+        constexpr index_t Ho = OutputSizes{}[0];
+        constexpr index_t Wo = OutputSizes{}[1];
+
+        constexpr index_t Y = FilterSizes{}[0];
+        constexpr index_t X = FilterSizes{}[1];
+
+        constexpr index_t ConvStrideH = ConvStrides{}[0];
+        constexpr index_t ConvStrideW = ConvStrides{}[1];
+
+        constexpr index_t ConvDilationH = ConvDilations{}[0];
+        constexpr index_t ConvDilationW = ConvDilations{}[1];
+
+        constexpr index_t E = C * Y * X;
+        constexpr index_t B = N * Ho * Wo;
+
+        // sanity-check for vectorized memory load
+        static_assert((Wo == 1 || (ConvStrideW == 1 || BlockCopyDataPerAccess_B == 1)) &&
+                          (X == 1 || ConvDilationW % BlockCopyDataPerAccess_B == 0),
+                      "wrong! aligment requirement for vectorized global load of input tensor will "
+                      "be violated");
+
+        // divide block work by [E, B]
+        static_assert(E % EPerBlock == 0 && B % BPerBlock == 0,
+                      "wrong! cannot divide work evenly among block");
+
+        constexpr index_t EBlockWork = E / EPerBlock;
+        constexpr index_t BBlockWork = B / BPerBlock;
+
+        constexpr auto block_work_desc =
+            make_cluster_descriptor(Sequence<EBlockWork, BBlockWork>{});
+
+        const auto block_work_id = block_work_desc.CalculateClusterIndex(get_block_1d_id());
+
+        const index_t e_block_data_on_global = block_work_id[Number<0>{}] * EPerBlock;
+        const index_t b_block_data_on_global = block_work_id[Number<1>{}] * BPerBlock;
+
+        // construct img_eb_global_desc
+        constexpr auto img_n_c_hip_wip_global_desc = transform_tensor_descriptor(
+            img_n_c_hi_wi_global_desc,
+            make_tuple(
+                PassThrough<N>{}, PassThrough<C>{}, Pad<Sequence<Hi, Wi>, LeftPads, RightPads>{}),
+            make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2, 3>{}),
+            make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2, 3>{}));
+
+        constexpr index_t Hip = img_n_c_hip_wip_global_desc.GetLengths()[2];
+        constexpr index_t Wip = img_n_c_hip_wip_global_desc.GetLengths()[3];
+
+        constexpr auto img_n_c_y_ho_x_wo_global_desc = transform_tensor_descriptor(
+            img_n_c_hip_wip_global_desc,
+            make_tuple(PassThrough<N>{},
+                       PassThrough<C>{},
+                       Embed<Hip, Sequence<Y, Ho>, Sequence<ConvDilationH, ConvStrideH, 0>>{},
+                       Embed<Wip, Sequence<X, Wo>, Sequence<ConvDilationW, ConvStrideW, 0>>{}),
+            make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}),
+            make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2, 3>{}, Sequence<4, 5>{}));
+
+        constexpr auto img_e_b_global_desc = transform_tensor_descriptor(
+            img_n_c_y_ho_x_wo_global_desc,
+            make_tuple(Merge<Sequence<C, Y, X>>{}, Merge<Sequence<N, Ho, Wo>>{}),
+            make_tuple(Sequence<1, 2, 4>{}, Sequence<0, 3, 5>{}),
+            make_tuple(Sequence<0>{}, Sequence<1>{}));
+
+        // blockwise atomic accumulation
+        auto blockwise_copy = BlockwiseGenericTensorSliceCopy_v4<BlockSize,
+                                                                 decltype(col_e_b_global_desc),
+                                                                 decltype(img_e_b_global_desc),
+                                                                 Sequence<EPerBlock, BPerBlock>,
+                                                                 BlockCopySubLengths_E_B,
+                                                                 BlockCopyClusterLengths_E_B,
+                                                                 BlockCopyThreadClusterArrangeOrder,
+                                                                 BlockCopySrcAccessOrder,
+                                                                 BlockCopyDstAccessOrder,
+                                                                 1,
+                                                                 1,
+                                                                 BlockCopyDataPerAccess_B,
+                                                                 BlockCopyDataPerAccess_B,
+                                                                 AddressSpace::Vgpr,
+                                                                 AddressSpace::Vgpr,
+                                                                 AddressSpace::Global,
+                                                                 InMemoryDataOperation::AtomicAdd>(
+            make_multi_index(e_block_data_on_global, b_block_data_on_global),
+            make_multi_index(e_block_data_on_global, b_block_data_on_global));
+
+        // blockwise copy
+        blockwise_copy.Run(p_col_global, p_img_global);
+    }
+};
+
+} // namespace ck
+#endif

driver/CMakeLists.txt

@@ -17,16 +17,17 @@ install(TARGETS host LIBRARY DESTINATION lib)
 if(DEVICE_BACKEND STREQUAL "AMD")
     set(CONV_SOURCE src/conv_driver.cpp)
     set(CONV_BWD_DATA_SOURCE src/conv_bwd_data_driver.cpp)
-    set(TRY_SOURCE src/try.cpp)
+    set(TRY_SOURCE src/col2im_driver.cpp)
 elseif(DEVICE_BACKEND STREQUAL "NVIDIA")
     set(CONV_SOURCE src/conv_driver.cu)
     set(CONV_BWD_DATA_SOURCE src/conv_bwd_data_driver.cu)
+    set(TRY_SOURCE src/col2im_driver.cu)
 endif()
 
 add_executable(conv_driver ${CONV_SOURCE}) 
 add_executable(conv_bwd_data_driver ${CONV_BWD_DATA_SOURCE}) 
-add_executable(try ${TRY_SOURCE}) 
+add_executable(col2im_driver ${TRY_SOURCE}) 
 
 target_link_libraries(conv_driver PRIVATE host)
 target_link_libraries(conv_bwd_data_driver PRIVATE host)
-target_link_libraries(try PRIVATE host)
+target_link_libraries(col2im_driver PRIVATE host)

driver/include/device_col2im_eb_nchw.hpp

+#pragma once
+#include <unistd.h>
+#include "device.hpp"
+#include "host_tensor.hpp"
+#include "gridwise_operation_wrapper.hpp"
+#include "gridwise_col2im_eb_nchw.hpp"
+
+template <typename T,
+          typename ColDesc,
+          typename ImgDesc,
+          typename FilterSizes,
+          typename OutputSizes,
+          typename ConvStrides,
+          typename ConvDilations,
+          typename LeftPads,
+          typename RightPads>
+void device_col2im_eb_nchw(ColDesc,
+                           const Tensor<T>& col_eb,
+                           ImgDesc,
+                           Tensor<T>& img_nchw,
+                           FilterSizes,
+                           OutputSizes,
+                           ConvStrides,
+                           ConvDilations,
+                           LeftPads,
+                           RightPads,
+                           std::size_t nrepeat)
+{
+    using namespace ck;
+
+    constexpr auto col_eb_desc   = ColDesc{};
+    constexpr auto img_nchw_desc = ImgDesc{};
+
+    constexpr index_t N  = img_nchw_desc.GetLengths()[0];
+    constexpr index_t C  = img_nchw_desc.GetLengths()[1];
+    constexpr index_t Hi = img_nchw_desc.GetLengths()[2];
+    constexpr index_t Wi = img_nchw_desc.GetLengths()[3];
+
+    constexpr index_t E = col_eb_desc.GetLengths()[0];
+    constexpr index_t B = col_eb_desc.GetLengths()[1];
+
+    std::size_t data_sz = sizeof(T);
+    DeviceMem col_eb_device_buf(data_sz * col_eb.mDesc.GetElementSpace());
+    DeviceMem img_nchw_device_buf(data_sz * img_nchw.mDesc.GetElementSpace());
+
+    col_eb_device_buf.ToDevice(col_eb.mData.data());
+    img_nchw_device_buf.ToDevice(img_nchw.mData.data());
+
+#if 1
+    constexpr index_t BlockSize = 256;
+
+    constexpr index_t EPerBlock = 128;
+    constexpr index_t BPerBlock = 128;
+
+    using BlockCopySubLengths_E_B            = Sequence<8, 8>;
+    using BlockCopyClusterLengths_E_B        = Sequence<16, 16>;
+    using BlockCopyThreadClusterArrangeOrder = Sequence<0, 1>; // [E, B]
+    using BlockCopySrcAccessOrder            = Sequence<0, 1>; // [E, B]
+    using BlockCopyDstAccessOrder            = Sequence<0, 1>; // [E, B]
+
+    constexpr index_t BlockCopyDataPerAccess_B = 1;
+#endif
+
+    constexpr index_t GridSize =
+        ((E + EPerBlock - 1) / EPerBlock) * ((B + BPerBlock - 1) / BPerBlock);
+
+    printf("%s: BlockSize %u, GridSize %u \n", __func__, BlockSize, GridSize);
+
+    constexpr auto gridwise_col2im = GridwiseCol2Im_eb_nchw<GridSize,
+                                                            BlockSize,
+                                                            T,
+                                                            ColDesc,
+                                                            ImgDesc,
+                                                            FilterSizes,
+                                                            OutputSizes,
+                                                            ConvStrides,
+                                                            ConvDilations,
+                                                            LeftPads,
+                                                            RightPads,
+                                                            EPerBlock,
+                                                            BPerBlock,
+                                                            BlockCopySubLengths_E_B,
+                                                            BlockCopyClusterLengths_E_B,
+                                                            BlockCopyThreadClusterArrangeOrder,
+                                                            BlockCopySrcAccessOrder,
+                                                            BlockCopyDstAccessOrder,
+                                                            BlockCopyDataPerAccess_B>{};
+
+    for(index_t i = 0; i < 1; ++i)
+    {
+        std::cout << "Start running " << nrepeat << " times..." << std::endl;
+
+        KernelTimer timer;
+        timer.Start();
+
+        for(index_t j = 0; j < nrepeat; ++j)
+        {
+            launch_kernel(run_gridwise_operation<decltype(gridwise_col2im),
+                                                 const T* const __restrict__,
+                                                 T* const __restrict__>,
+                          dim3(GridSize),
+                          dim3(BlockSize),
+                          0,
+                          0,
+                          const_cast<const T* const __restrict__>(
+                              static_cast<T*>(col_eb_device_buf.GetDeviceBuffer())),
+                          const_cast<T* const __restrict__>(
+                              static_cast<T*>(img_nchw_device_buf.GetDeviceBuffer())));
+        }
+
+        timer.End();
+
+        float ave_time = timer.GetElapsedTime() / nrepeat;
+
+        std::cout << "Average time : " << ave_time << " ms" << std::endl;
+    }
+
+    img_nchw_device_buf.FromDevice(img_nchw.mData.data());
+}

driver/include/host_col2im.hpp

+#pragma once
+#include "host_tensor.hpp"
+
+template <typename T,
+          typename FilterSizes,
+          typename OutputSizes,
+          typename ConvStrides,
+          typename ConvDilations,
+          typename LeftPads,
+          typename RightPads>
+void host_col2im(const Tensor<T>& in_eb,
+                 Tensor<T>& in_nchw,
+                 FilterSizes,
+                 OutputSizes,
+                 ConvStrides,
+                 ConvDilations,
+                 LeftPads,
+                 RightPads)
+{
+    using namespace ck;
+
+    int N  = in_nchw.mDesc.GetLengths()[0];
+    int C  = in_nchw.mDesc.GetLengths()[1];
+    int HI = in_nchw.mDesc.GetLengths()[2];
+    int WI = in_nchw.mDesc.GetLengths()[3];
+
+    int Y = FilterSizes{}[0];
+    int X = FilterSizes{}[1];
+
+    int HO = OutputSizes{}[0];
+    int WO = OutputSizes{}[1];
+
+    auto f = [&](auto n, auto c, auto hi, auto wi) {
+        double v = 0;
+
+        for(int y = 0; y < Y; ++y)
+        {
+            int h_tmp = hi + LeftPads{}[0] - y * ConvDilations{}[0];
+
+            if(h_tmp >= 0 && h_tmp < HI && h_tmp % ConvStrides{}[0] == 0)
+            {
+                int ho = h_tmp / ConvStrides{}[0];
+
+                for(int x = 0; x < X; ++x)
+                {
+                    int w_tmp = wi + LeftPads{}[1] - x * ConvDilations{}[1];
+
+                    if(w_tmp >= 0 && w_tmp < WI && w_tmp % ConvStrides{}[1] == 0)
+                    {
+                        int wo = w_tmp / ConvStrides{}[1];
+
+                        int e = c * (Y * X) + y * X + x;
+                        int b = n * (HO * WO) + ho * WO + wo;
+
+                        v += in_eb(e, b);
+                    }
+                }
+            }
+        }
+
+        in_nchw(n, c, hi, wi) = v;
+    };
+
+    auto f_par = make_ParallelTensorFunctor(f,
+                                            in_nchw.mDesc.GetLengths()[0],
+                                            in_nchw.mDesc.GetLengths()[1],
+                                            in_nchw.mDesc.GetLengths()[2],
+                                            in_nchw.mDesc.GetLengths()[3]);
+
+    f_par(std::thread::hardware_concurrency());
+}

driver/src/col2im_driver.cpp

+#include <iostream>
+#include <numeric>
+#include <initializer_list>
+#include <cstdlib>
+#include <stdlib.h>
+#include "config.hpp"
+#include "print.hpp"
+#include "device.hpp"
+#include "host_tensor_generator.hpp"
+#include "conv_common.hpp"
+#include "host_conv.hpp"
+#include "device_tensor.hpp"
+#include "host_col2im.hpp"
+#include "device_col2im_eb_nchw.hpp"
+
+int main(int argc, char* argv[])
+{
+    using namespace ck;
+
+#if 1
+    constexpr index_t N  = 2;
+    constexpr index_t C  = 8;
+    constexpr index_t HI = 8;
+    constexpr index_t WI = 8;
+    constexpr index_t K  = 128;
+    constexpr index_t Y  = 4;
+    constexpr index_t X  = 4;
+
+    using ConvStrides   = Sequence<1, 1>;
+    using ConvDilations = Sequence<1, 1>;
+
+    using LeftPads  = Sequence<1, 1>;
+    using RightPads = Sequence<2, 2>;
+#elif 0
+    // 3x3, 34x34
+    constexpr index_t N  = 64;
+    constexpr index_t C  = 256;
+    constexpr index_t HI = 34;
+    constexpr index_t WI = 34;
+    constexpr index_t K  = 128;
+    constexpr index_t Y  = 3;
+    constexpr index_t X  = 3;
+
+    using ConvStrides   = Sequence<1, 1>;
+    using ConvDilations = Sequence<1, 1>;
+
+    using LeftPads  = Sequence<0, 0>;
+    using RightPads = Sequence<0, 0>;
+#elif 0
+    // 1x1 filter, 8x8 image
+    // cudnn@V100 68%, ck@V100 72%, ck@P100 52%, ck@VII 42%
+    constexpr index_t N  = 64;
+    constexpr index_t C  = 1536;
+    constexpr index_t HI = 8;
+    constexpr index_t WI = 8;
+    constexpr index_t K  = 256;
+    constexpr index_t Y  = 1;
+    constexpr index_t X  = 1;
+
+    using ConvStrides   = Sequence<1, 1>;
+    using ConvDilations = Sequence<1, 1>;
+
+    using LeftPads  = Sequence<0, 0>;
+    using RightPads = Sequence<0, 0>;
+#elif 0
+    // 1x1 filter, 8x8 image
+    // cudnn@V100 77%, ck@V100 76%, ck@P100 79%, ck@VII 51%
+    constexpr index_t N  = 128;
+    constexpr index_t C  = 2048;
+    constexpr index_t HI = 8;
+    constexpr index_t WI = 8;
+    constexpr index_t K  = 384;
+    constexpr index_t Y  = 1;
+    constexpr index_t X  = 1;
+
+    using ConvStrides   = Sequence<1, 1>;
+    using ConvDilations = Sequence<1, 1>;
+
+    using LeftPads  = Sequence<0, 0>;
+    using RightPads = Sequence<0, 0>;
+#elif 0
+    // 1x1 filter, 7x7 image
+    // cudnn@V100 82%, ck@V100 76%, ck@P100 67%, ck@VII 64%
+    constexpr index_t N  = 128;
+    constexpr index_t C  = 832;
+    constexpr index_t HI = 7;
+    constexpr index_t WI = 7;
+    constexpr index_t K  = 384;
+    constexpr index_t Y  = 1;
+    constexpr index_t X  = 1;
+
+    using ConvStrides   = Sequence<1, 1>;
+    using ConvDilations = Sequence<1, 1>;
+
+    using LeftPads  = Sequence<0, 0>;
+    using RightPads = Sequence<0, 0>;
+#elif 0
+    // 1x1 filter, 8x8 image
+    // cudnn@V100 83%, ck@V100 75%, ck@P100 78%, ck@VII 65%
+    constexpr index_t N  = 128;
+    constexpr index_t C  = 1280;
+    constexpr index_t HI = 8;
+    constexpr index_t WI = 8;
+    constexpr index_t K  = 384;
+    constexpr index_t Y  = 1;
+    constexpr index_t X  = 1;
+
+    using ConvStrides   = Sequence<1, 1>;
+    using ConvDilations = Sequence<1, 1>;
+
+    using LeftPads  = Sequence<0, 0>;
+    using RightPads = Sequence<0, 0>;
+#elif 0
+    // 1x1 filter, 14x14 image
+    // cudnn@V100 62%, ck@V100 68%, ck@P100 70%, ck@VII 50%
+    constexpr index_t N  = 128;
+    constexpr index_t C  = 512;
+    constexpr index_t HI = 14;
+    constexpr index_t WI = 14;
+    constexpr index_t K  = 128;
+    constexpr index_t Y  = 1;
+    constexpr index_t X  = 1;
+
+    using ConvStrides   = Sequence<1, 1>;
+    using ConvDilations = Sequence<1, 1>;
+
+    using LeftPads  = Sequence<0, 0>;
+    using RightPads = Sequence<0, 0>;
+#elif 0
+    // 1x1 filter, 8x8 image
+    // cudnn@V100 74%, ck@V100 57%, ck@P100 78%, ck@VII 61%
+    constexpr index_t N  = 64;
+    constexpr index_t C  = 1536;
+    constexpr index_t HI = 8;
+    constexpr index_t WI = 8;
+    constexpr index_t K  = 384;
+    constexpr index_t Y  = 1;
+    constexpr index_t X  = 1;
+
+    using ConvStrides   = Sequence<1, 1>;
+    using ConvDilations = Sequence<1, 1>;
+
+    using LeftPads  = Sequence<0, 0>;
+    using RightPads = Sequence<0, 0>;
+#elif 0
+    // 1x1 filter, 28x28 image
+    // cudnn@V100 86%, ck@V100 84%, ck@P100 80%, ck@VII 69%
+    constexpr index_t N  = 128;
+    constexpr index_t C  = 256;
+    constexpr index_t HI = 28;
+    constexpr index_t WI = 28;
+    constexpr index_t K  = 128;
+    constexpr index_t Y  = 1;
+    constexpr index_t X  = 1;
+
+    using ConvStrides   = Sequence<1, 1>;
+    using ConvDilations = Sequence<1, 1>;
+
+    using LeftPads  = Sequence<0, 0>;
+    using RightPads = Sequence<0, 0>;
+#elif 0
+    // 1x1 filter, 7x7 image
+    // cudnn@V100 71%, ck@V100 55%, ck@P100 70%, ck@VII 62%
+    constexpr index_t N  = 128;
+    constexpr index_t C  = 832;
+    constexpr index_t HI = 7;
+    constexpr index_t WI = 7;
+    constexpr index_t K  = 256;
+    constexpr index_t Y  = 1;
+    constexpr index_t X  = 1;
+
+    using ConvStrides   = Sequence<1, 1>;
+    using ConvDilations = Sequence<1, 1>;
+
+    using LeftPads  = Sequence<0, 0>;
+    using RightPads = Sequence<0, 0>;
+#elif 0
+    // 1x1 filter, 17x17 input
+    // cudnn@V100 81%, ck@V100 76%, ck@P100 70%, ck@VII 76%
+    constexpr index_t N  = 128;
+    constexpr index_t C  = 768;
+    constexpr index_t HI = 17;
+    constexpr index_t WI = 17;
+    constexpr index_t K  = 128;
+    constexpr index_t Y  = 1;
+    constexpr index_t X  = 1;
+
+    using ConvStrides   = Sequence<1, 1>;
+    using ConvDilations = Sequence<1, 1>;
+
+    using LeftPads  = Sequence<0, 0>;
+    using RightPads = Sequence<0, 0>;
+#elif 0
+    // 1x1 filter, 14x14 image
+    // cudnn@V100 73%, ck@V100 71%, ck@P100 70%, ck@VII 64%
+    constexpr index_t N  = 128;
+    constexpr index_t C  = 528;
+    constexpr index_t HI = 14;
+    constexpr index_t WI = 14;
+    constexpr index_t K  = 128;
+    constexpr index_t Y  = 1;
+    constexpr index_t X  = 1;
+
+    using ConvStrides   = Sequence<1, 1>;
+    using ConvDilations = Sequence<1, 1>;
+
+    using LeftPads  = Sequence<0, 0>;
+    using RightPads = Sequence<0, 0>;
+#elif 0
+    // 1x1 filter, 14x14 image
+    // cudnn@V100 73%, ck@V100 72%, ck@P100 79%, ck@VII 75%
+    constexpr index_t N  = 128;
+    constexpr index_t C  = 528;
+    constexpr index_t HI = 14;
+    constexpr index_t WI = 14;
+    constexpr index_t K  = 256;
+    constexpr index_t Y  = 1;
+    constexpr index_t X  = 1;
+
+    using ConvStrides   = Sequence<1, 1>;
+    using ConvDilations = Sequence<1, 1>;
+
+    using LeftPads  = Sequence<0, 0>;
+    using RightPads = Sequence<0, 0>;
+#elif 0
+    // 1x1 filter, 7x7 image
+    // cudnn@V100 49%, ck@V100 50%, ck@P100 61%, ck@VII 52%
+    constexpr index_t N  = 128;
+    constexpr index_t C  = 832;
+    constexpr index_t HI = 7;
+    constexpr index_t WI = 7;
+    constexpr index_t K  = 128;
+    constexpr index_t Y  = 1;
+    constexpr index_t X  = 1;
+
+    using ConvStrides   = Sequence<1, 1>;
+    using ConvDilations = Sequence<1, 1>;
+
+    using LeftPads  = Sequence<0, 0>;
+    using RightPads = Sequence<0, 0>;
+#elif 0
+    // 3x3 filter, 2x2 stride, 35x35 input, 17x17 output
+    // cudnn@V100 90%, ck@V100 93%, ck@P100 83%, ck@VII 81%
+    constexpr index_t N  = 128;
+    constexpr index_t C  = 288;
+    constexpr index_t HI = 35;
+    constexpr index_t WI = 35;
+    constexpr index_t K  = 384;
+    constexpr index_t Y  = 3;
+    constexpr index_t X  = 3;
+
+    using ConvStrides   = Sequence<2, 2>;
+    using ConvDilations = Sequence<1, 1>;
+
+    using LeftPads  = Sequence<0, 0>;
+    using RightPads = Sequence<0, 0>;
+#elif 0
+    // 5x5 filter, 2x2 pad, 7x7 input
+    constexpr index_t N  = 128;
+    constexpr index_t C  = 48;
+    constexpr index_t HI = 7;
+    constexpr index_t WI = 7;
+    constexpr index_t K  = 128;
+    constexpr index_t Y  = 5;
+    constexpr index_t X  = 5;
+
+    using ConvStrides   = Sequence<1, 1>;
+    using ConvDilations = Sequence<1, 1>;
+
+    using LeftPads  = Sequence<2, 2>;
+    using RightPads = Sequence<2, 2>;
+#elif 0
+    // 7x1 filter, 3x0 pad, 17x17 input
+    constexpr index_t N  = 128;
+    constexpr index_t C  = 128;
+    constexpr index_t HI = 17;
+    constexpr index_t WI = 17;
+    constexpr index_t K  = 128;
+    constexpr index_t Y  = 7;
+    constexpr index_t X  = 1;
+
+    using ConvStrides   = Sequence<1, 1>;
+    using ConvDilations = Sequence<1, 1>;
+
+    using LeftPads  = Sequence<3, 0>;
+    using RightPads = Sequence<3, 0>;
+#elif 1
+    // 1x7 filter, 0x3 pad, 17x17 input
+    constexpr index_t N  = 128;
+    constexpr index_t C  = 128;
+    constexpr index_t HI = 17;
+    constexpr index_t WI = 17;
+    constexpr index_t K  = 128;
+    constexpr index_t Y  = 1;
+    constexpr index_t X  = 7;
+
+    using ConvStrides   = Sequence<1, 1>;
+    using ConvDilations = Sequence<1, 1>;
+
+    using LeftPads  = Sequence<0, 3>;
+    using RightPads = Sequence<0, 3>;
+#endif
+
+    constexpr auto img_nchw_desc = make_native_tensor_descriptor_packed(Sequence<N, C, HI, WI>{});
+    constexpr auto wei_kcyx_desc = make_native_tensor_descriptor_packed(Sequence<K, C, Y, X>{});
+    constexpr auto out_nkhw_desc = get_convolution_output_default_4d_tensor_descriptor(
+        img_nchw_desc, wei_kcyx_desc, ConvStrides{}, ConvDilations{}, LeftPads{}, RightPads{});
+
+    constexpr index_t HO = out_nkhw_desc.GetLengths()[2];
+    constexpr index_t WO = out_nkhw_desc.GetLengths()[3];
+
+    constexpr auto col_eb_desc =
+        make_native_tensor_descriptor_packed(Sequence<C * Y * X, N * HO * WO>{});
+
+    using FilterSizes = Sequence<Y, X>;
+    using OutputSizes = Sequence<HO, WO>;
+
+    ostream_tensor_descriptor(col_eb_desc, std::cout << "col_eb_desc: ");
+    ostream_tensor_descriptor(img_nchw_desc, std::cout << "img_nchw_desc: ");
+    print_array("FilterSizes", FilterSizes{});
+    print_array("OutputSizes", OutputSizes{});
+    print_array("LeftPads", LeftPads{});
+    print_array("LeftPads", LeftPads{});
+    print_array("RightPads", RightPads{});
+    print_array("ConvStrides", ConvStrides{});
+    print_array("ConvDilations", ConvDilations{});
+
+    Tensor<float> col_eb(make_HostTensorDescriptor(col_eb_desc));
+    Tensor<float> img_nchw_host(make_HostTensorDescriptor(img_nchw_desc));
+    Tensor<float> img_nchw_device(make_HostTensorDescriptor(img_nchw_desc));
+
+    std::size_t num_thread = std::thread::hardware_concurrency();
+
+    if(argc != 3)
+    {
+        printf("arg1: do_verification, arg2: nrepeat\n");
+        exit(1);
+    }
+
+    bool do_verification = atoi(argv[1]);
+    std::size_t nrepeat  = atoi(argv[2]);
+
+    if(do_verification)
+    {
+#if 0
+        col_eb.GenerateTensorValue(GeneratorTensor_1{}, num_thread);
+#else
+        col_eb.GenerateTensorValue(GeneratorTensor_2{-5, 5}, num_thread);
+#endif
+    }
+
+    device_col2im_eb_nchw(col_eb_desc,
+                          col_eb,
+                          img_nchw_desc,
+                          img_nchw_device,
+                          FilterSizes{},
+                          OutputSizes{},
+                          ConvStrides{},
+                          ConvDilations{},
+                          LeftPads{},
+                          RightPads{},
+                          nrepeat);
+
+    if(do_verification)
+    {
+        host_col2im(col_eb,
+                    img_nchw_host,
+                    FilterSizes{},
+                    OutputSizes{},
+                    ConvStrides{},
+                    ConvDilations{},
+                    LeftPads{},
+                    RightPads{});
+
+        check_error(img_nchw_host, img_nchw_device);
+
+#if 0
+        LogRange(std::cout << "col_eb : ", col_eb.mData, ",") << std::endl;
+        LogRange(std::cout << "img_nchw_host : ", img_nchw_host.mData, ",") << std::endl;
+        LogRange(std::cout << "img_nchw_device : ", img_nchw_device.mData, ",") << std::endl;
+#endif
+    }
+}

driver/src/try.cpp

-#include <iostream>
-#include <numeric>
-#include <initializer_list>
-#include <cstdlib>
-#include <stdlib.h>
-#include <half.hpp>
-#include "config.hpp"
-#include "print.hpp"
-#include "device.hpp"
-#include "host_tensor_generator.hpp"
-#include "conv_common.hpp"
-#include "host_conv.hpp"
-#include "device_tensor.hpp"
-#include "device_convolution_implicit_gemm_v4r1_nchw_kcyx_nkhw.hpp"
-#include "device_convolution_implicit_gemm_v4r4_nchw_kcyx_nkhw.hpp"
-
-int main(int argc, char* argv[])
-{
-    using namespace ck;
-
-    auto idx1 = std::array<index_t, 2>{{1, 0}};
-    auto idx2 = Array<index_t, 2>{{1, 0}};
-    auto idx3 = MultiIndex<2>{{1, 0}};
-
-    auto idx0 = MultiIndex<2>{{1, 0}};
-
-    print_array("idx2", idx2);
-    print_array("idx3", idx2);
-}