Add image to column kernel

example/52_image_to_column/CMakeLists.txt

+list(APPEND gpu_list gfx908 gfx90a gfx940 gfx941 gfx942)
+set(target 0)
+foreach(gpu IN LISTS GPU_TARGETS)
+ if(gpu IN_LIST gpu_list AND target EQUAL 0)
+   add_custom_target(example_image_to_column)
+   add_example_executable(example_image_to_column_f32 image_to_column_f32.cpp)
+   add_dependencies(example_image_to_column example_image_to_column_f32)
+   set(target 1)
+ endif()
+endforeach()

example/52_image_to_column/common.hpp

+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
+
+#pragma once
+
+#include <cstdlib>
+#include <initializer_list>
+#include <iostream>
+#include <numeric>
+
+#include "ck/ck.hpp"
+#include "ck/tensor_operation/gpu/device/impl/device_image_to_column_impl.hpp"
+#include "ck/tensor_operation/gpu/device/tensor_layout.hpp"
+
+#include "ck/library/utility/algorithm.hpp"
+#include "ck/library/utility/check_err.hpp"
+#include "ck/library/utility/convolution_host_tensor_descriptor_helper.hpp"
+#include "ck/library/utility/convolution_parameter.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/reference_tensor_operation/cpu/reference_image_to_column.hpp"
+
+template <ck::index_t... Is>
+using S = ck::Sequence<Is...>;
+
+static inline constexpr ck::index_t NDimSpatial = 2;
+
+using FP32 = float;
+
+struct ExecutionConfig final
+{
+    bool do_verification = true;
+    int init_method      = 1;
+    bool time_kernel     = true;
+};
+
+#define DefaultConvParams                                                            \
+    ck::utils::conv::ConvParam                                                       \
+    {                                                                                \
+        NDimSpatial, 1, 32, 1, 1, {4, 4}, {64, 64}, {1, 1}, {1, 1}, {0, 0}, { 0, 0 } \
+    }
+
+inline void print_help_msg()
+{
+    std::cerr << "arg1: verification (0=no, 1=yes)\n"
+              << "arg2: initialization (0=no init, 1=integer value, 2=decimal value)\n"
+              << "arg3: time kernel (0=no, 1=yes)\n"
+              << ck::utils::conv::get_conv_param_parser_helper_msg() << std::endl;
+}
+
+inline bool parse_cmd_args(int argc,
+                           char* argv[],
+                           ExecutionConfig& config,
+                           ck::utils::conv::ConvParam& conv_params)
+{
+    constexpr int num_execution_config_args =
+        3; // arguments for do_verification, init_method, time_kernel
+    constexpr int num_conv_param_leading_args = 5; // arguments for num_dim_spatial_, G_, N_, K_, C_
+
+    constexpr int threshold_to_catch_partial_args = 1 + num_execution_config_args;
+    constexpr int threshold_to_catch_all_args =
+        threshold_to_catch_partial_args + num_conv_param_leading_args;
+
+    if(argc == 1)
+    {
+        // use default
+        config = ExecutionConfig{};
+    }
+    // catch only ExecutionConfig arguments
+    else if(argc == threshold_to_catch_partial_args)
+    {
+        config.do_verification = std::stoi(argv[1]);
+        config.init_method     = std::stoi(argv[2]);
+        config.time_kernel     = std::stoi(argv[3]);
+    }
+    // catch both ExecutionConfig & ConvParam arguments
+    else if(threshold_to_catch_all_args < argc && ((argc - threshold_to_catch_all_args) % 3 == 0))
+    {
+        config.do_verification = std::stoi(argv[1]);
+        config.init_method     = std::stoi(argv[2]);
+        config.time_kernel     = std::stoi(argv[3]);
+
+        const ck::index_t num_dim_spatial = std::stoi(argv[4]);
+        conv_params                       = ck::utils::conv::parse_conv_param(
+            num_dim_spatial, threshold_to_catch_partial_args, argv);
+    }
+    else
+    {
+        print_help_msg();
+        return false;
+    }
+
+    return true;
+}

example/52_image_to_column/image_to_column_f32.cpp

+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
+
+#include "common.hpp"
+
+using InDataType  = FP32;
+using OutDataType = FP32;
+
+using InLayout = ck::tensor_layout::convolution::GNHWC;
+
+// clang-format off
+using DeviceImgToColInstance = ck::tensor_operation::device::DeviceImageToColumnImpl
+// ######| NDimSpatial|   ALayout|   BLayout|   DsLayout|  ELayout|       AData|       BData|     AccData|         CShuffle|       DsData|      EData| AElementwise| BElementwise| CDEElementwise| ConvolutionBackward| DoPad| DoPad|      NumGemmK| Block|  MPer|  NPer|  KPer| AK1| BK1| MPer| NPer|    MXdl|    NXdl|    ABlockTransfer| ABlockTransfer| ABlockTransfer| ABlockTransfer| ABlockTransfer| ABlockTransfer| ABlockLds|    BBlockTransfer| BBlockTransfer| BBlockTransfer| BBlockTransfer| BBlockTransfer| BBlockTransfer| BBlockLds| CShuffleMXdl| CShuffleNXdl|   CDEBlockTransfer| CDEBlockTransfer|
+// ######|            |          |          |           |         |        Type|        Type|        Type|         DataType|         Type|       Type|    Operation|    Operation|      Operation|  DataSpecialization| GemmM| GemmN| PrefetchStage|  Size| Block| Block| Block|    |    |  XDL|  XDL| PerWave| PerWave|     ThreadCluster|  ThreadCluster| SrcAccessOrder|   SrcVectorDim|      SrcScalar|      DstScalar|    ExtraM|     ThreadCluster|  ThreadCluster| SrcAccessOrder|   SrcVectorDim|      SrcScalar|      DstScalar|    ExtraN|      PerWave|      PerWave|  _MBlock_MPerBlock|  ScalarPerVector|
+// ######|            |          |          |           |         |            |            |            |                 |             |           |             |             |               |                    |      |      |              |      |      |      |      |    |    |     |     |        |        | Lengths_AK0_M_AK1|   ArrangeOrder|               |               |      PerVector|  PerVector_AK1|          | Lengths_BK0_N_BK1|   ArrangeOrder|               |               |      PerVector|  PerVector_BK1|          |   PerShuffle|   PerShuffle|  _NBlock_NPerBlock|       _NPerBlock|
+// ######|            |          |          |           |         |            |            |            |                 |             |           |             |             |               |                    |      |      |              |      |      |      |      |    |    |     |     |        |        |                  |               |               |               |               |               |          |                  |               |               |               |               |               |          |             |             |                   |                 |
+         < NDimSpatial, InLayout, InDataType, OutDataType, 256, 128, 128, S<128, 128>, S<16, 16>, 4>;
+// clang-format on
+
+bool RunImageToColumn(const ExecutionConfig& config, const ck::utils::conv::ConvParam& conv_params)
+{
+
+    const auto N = conv_params.N_;
+    const auto C = conv_params.C_;
+
+    const ck::index_t NDoHoWo =
+        N * ck::accumulate_n<ck::index_t>(
+                conv_params.output_spatial_lengths_.begin(), NDimSpatial, 1, std::multiplies<>());
+    const ck::index_t CZYX =
+        C * ck::accumulate_n<ck::index_t>(
+                conv_params.filter_spatial_lengths_.begin(), NDimSpatial, 1, std::multiplies<>());
+
+    const auto in_desc =
+        ck::utils::conv::make_input_host_tensor_descriptor_g_n_c_wis_packed<InLayout>(conv_params);
+    const auto out_desc = HostTensorDescriptor({NDoHoWo, CZYX});
+
+    std::array<ck::index_t, NDimSpatial> input_spatial_lengths{};
+    std::array<ck::index_t, NDimSpatial> filter_spatial_lengths{};
+    std::array<ck::index_t, NDimSpatial> output_spatial_lengths{};
+    std::array<ck::index_t, NDimSpatial + 3> input_g_n_c_wis_strides{};
+    std::array<ck::index_t, 2> output_m_k_strides{};
+    std::array<ck::index_t, NDimSpatial> conv_filter_strides{};
+    std::array<ck::index_t, NDimSpatial> conv_filter_dilations{};
+    std::array<ck::index_t, NDimSpatial> input_left_pads{};
+    std::array<ck::index_t, NDimSpatial> input_right_pads{};
+
+    auto copy = [](const auto& x, auto& y) { std::copy(x.begin(), x.end(), y.begin()); };
+
+    copy(conv_params.input_spatial_lengths_, input_spatial_lengths);
+    copy(conv_params.filter_spatial_lengths_, filter_spatial_lengths);
+    copy(conv_params.output_spatial_lengths_, output_spatial_lengths);
+    copy(in_desc.GetStrides(), input_g_n_c_wis_strides);
+    copy(out_desc.GetStrides(), output_m_k_strides);
+    copy(conv_params.conv_filter_strides_, conv_filter_strides);
+    copy(conv_params.conv_filter_dilations_, conv_filter_dilations);
+    copy(conv_params.input_left_pads_, input_left_pads);
+    copy(conv_params.input_right_pads_, input_right_pads);
+
+    Tensor<InDataType> in(in_desc);
+    Tensor<OutDataType> out_device(out_desc);
+    Tensor<OutDataType> out_host(out_desc);
+
+    std::cout << "in: " << in.mDesc << std::endl;
+    std::cout << "out: " << out_device.mDesc << std::endl;
+
+    switch(config.init_method)
+    {
+    case 0: break;
+    case 1: in.GenerateTensorValue(GeneratorTensor_2<InDataType>{-5, 5}); break;
+    default: in.GenerateTensorValue(GeneratorTensor_3<InDataType>{-0.5, 0.5});
+    }
+
+    DeviceMem in_device_buf(sizeof(InDataType) * in.mDesc.GetElementSpaceSize());
+    DeviceMem out_device_buf(sizeof(OutDataType) * out_device.mDesc.GetElementSpaceSize());
+
+    in_device_buf.ToDevice(in.mData.data());
+
+    // reset input to zero
+    out_device_buf.SetZero();
+
+    static_assert(std::is_default_constructible_v<DeviceImgToColInstance>);
+
+    // do conv
+    auto img2col  = DeviceImgToColInstance{};
+    auto invoker  = img2col.MakeInvoker();
+    auto argument = img2col.MakeArgument(in_device_buf.GetDeviceBuffer(),
+                                         out_device_buf.GetDeviceBuffer(),
+                                         N,
+                                         C,
+                                         input_spatial_lengths,
+                                         filter_spatial_lengths,
+                                         output_spatial_lengths,
+                                         input_g_n_c_wis_strides,
+                                         output_m_k_strides,
+                                         conv_filter_strides,
+                                         conv_filter_dilations,
+                                         input_left_pads,
+                                         input_right_pads);
+
+    if(!img2col.IsSupportedArgument(argument))
+    {
+        std::cerr << "wrong! device_img2col with the specified compilation parameters does "
+                     "not support this img2col problem"
+                  << std::endl;
+
+        return false;
+    }
+
+    float ave_time = invoker.Run(argument, StreamConfig{nullptr, config.time_kernel});
+
+    std::size_t num_btype = NDoHoWo * CZYX * sizeof(InDataType);
+
+    float gb_per_sec = num_btype / 1.E6 / ave_time;
+
+    std::cout << "Perf: " << ave_time << " ms, " << gb_per_sec << " GB/s" << std::endl;
+
+    if(config.do_verification)
+    {
+        auto ref_image_to_column = ck::tensor_operation::host::
+            ReferenceImageToColumn<NDimSpatial, InLayout, InDataType, OutDataType>();
+
+        auto ref_invoker = ref_image_to_column.MakeInvoker();
+
+        auto ref_argument = ref_image_to_column.MakeArgument(in,
+                                                             out_host,
+                                                             conv_params.filter_spatial_lengths_,
+                                                             conv_params.conv_filter_strides_,
+                                                             conv_params.conv_filter_dilations_,
+                                                             conv_params.input_left_pads_,
+                                                             conv_params.input_right_pads_);
+
+        if(!ref_image_to_column.IsSupportedArgument(&ref_argument))
+        {
+            std::cerr << "wrong! ref_img2col with the specified compilation parameters does "
+                         "not support this img2col problem"
+                      << std::endl;
+            return false;
+        }
+
+        ref_invoker.Run(ref_argument);
+
+        out_device_buf.FromDevice(out_device.mData.data());
+
+        return ck::utils::check_err(out_device.mData, out_host.mData);
+    }
+
+    return true;
+}
+
+int RunImageToColumnExample(int argc, char* argv[])
+{
+    ExecutionConfig config;
+    ck::utils::conv::ConvParam conv_params = DefaultConvParams;
+
+    if(!parse_cmd_args(argc, argv, config, conv_params))
+    {
+        return EXIT_FAILURE;
+    }
+
+    if(conv_params.num_dim_spatial_ != NDimSpatial)
+    {
+        std::cerr << "unsupported # of spatials dimensions" << std::endl;
+        return EXIT_FAILURE;
+    }
+
+    return !RunImageToColumn(config, conv_params);
+}
+
+int main(int argc, char* argv[]) { return RunImageToColumnExample(argc, argv); }

include/ck/tensor_operation/gpu/device/device_image_to_column.hpp

+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
+
+#pragma once
+
+#include <array>
+
+#include "ck/tensor_operation/gpu/device/device_base.hpp"
+
+namespace ck {
+namespace tensor_operation {
+namespace device {
+
+// Image to column:
+//   input : input image [N, Di, Hi, Wi, C],
+//   output : output image [N * Do * Ho * Wo, Z *  Y * X * C]
+template <index_t NDimSpatial,
+          typename InputLayout,
+          typename InputDataType,
+          typename OutputDataType>
+struct DeviceImageToColumn : public BaseOperator
+{
+    virtual std::unique_ptr<BaseArgument>
+    MakeArgumentPointer(const void* p_in, // input image
+                        void* p_out,      // output image
+                        const ck::index_t N,
+                        const ck::index_t C,
+                        const std::array<index_t, NDimSpatial>& input_spatial_lengths,
+                        const std::array<index_t, NDimSpatial>& filter_spatial_lengths,
+                        const std::array<index_t, NDimSpatial>& output_spatial_lengths,
+                        const std::array<index_t, NDimSpatial + 3>& input_g_n_c_wis_strides,
+                        const std::array<index_t, 2>& output_m_k_strides,
+                        const std::array<index_t, NDimSpatial>& conv_filter_strides,
+                        const std::array<index_t, NDimSpatial>& conv_filter_dilations,
+                        const std::array<index_t, NDimSpatial>& input_left_pads,
+                        const std::array<index_t, NDimSpatial>& input_right_pads) = 0;
+
+    virtual std::unique_ptr<BaseInvoker> MakeInvokerPointer() = 0;
+};
+
+} // namespace device
+} // namespace tensor_operation
+} // namespace ck

include/ck/tensor_operation/gpu/grid/gridwise_image_to_column.hpp

+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
+
+#pragma once
+
+#include "ck/utility/common_header.hpp"
+#include "ck/tensor_description/multi_index_transform_helper.hpp"
+#include "ck/tensor_description/tensor_descriptor.hpp"
+#include "ck/tensor_description/tensor_descriptor_helper.hpp"
+#include "ck/tensor_operation/gpu/grid/block_to_ctile_map.hpp"
+#include "ck/tensor_operation/gpu/grid/gridwise_gemm_pipeline_selector.hpp"
+#include "ck/tensor_operation/gpu/block/blockwise_gemm_xdlops.hpp"
+#include "ck/tensor_operation/gpu/block/thread_group_tensor_slice_transfer_v4r1.hpp"
+#include "ck/tensor_operation/gpu/block/thread_group_tensor_slice_transfer_v7.hpp"
+#include "ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer.hpp"
+#include "ck/tensor_operation/gpu/element/element_wise_operation.hpp"
+
+namespace ck {
+
+template <typename InputGridDesc,
+          typename InputDataType,
+          typename OutputGridDesc,
+          typename OutputDataType,
+          index_t BlockSize,
+          index_t MPerBlock,
+          index_t KPerBlock,
+          typename SliceLengths,
+          typename ThreadClusterLengths,
+          index_t ScalarPerVector,
+          typename Block2ETileMap>
+struct GridwiseImageToColumn
+{
+
+    static constexpr auto I0 = Number<0>{};
+    static constexpr auto I1 = Number<1>{};
+
+    using ThisThreadBlock = ThisThreadBlock<BlockSize>;
+
+    __device__ static void Run(const InputGridDesc& in_grid_desc,
+                               const InputDataType* __restrict__ p_in_global,
+                               const OutputGridDesc& out_grid_desc,
+                               OutputDataType* __restrict__ p_out_global,
+                               const Block2ETileMap& block_2_tile_map)
+    {
+        const auto block_work_idx =
+            block_2_tile_map.CalculateBottomIndex(make_multi_index(get_block_1d_id()));
+
+        const index_t m_block_data_idx_on_grid =
+            __builtin_amdgcn_readfirstlane(block_work_idx[I0] * MPerBlock);
+
+        const index_t k_block_data_idx_on_grid =
+            __builtin_amdgcn_readfirstlane(block_work_idx[I1] * KPerBlock);
+
+        // Global Memory
+        const auto in_global_buf = make_dynamic_buffer<AddressSpaceEnum::Global>(
+            p_in_global, in_grid_desc.GetElementSpaceSize());
+        auto out_global_buf = make_dynamic_buffer<AddressSpaceEnum::Global>(
+            p_out_global, out_grid_desc.GetElementSpaceSize());
+
+        auto copy_global_to_global = ThreadGroupTensorSliceTransfer_v7<
+            ThisThreadBlock,
+            Tuple<InputDataType>,
+            Tuple<OutputDataType>,
+            decltype(tie(in_grid_desc)),
+            decltype(tie(out_grid_desc)),
+            tensor_operation::element_wise::PassThrough,
+            Sequence<static_cast<index_t>(InMemoryDataOperationEnum::Set)>,
+            SliceLengths,
+            ThreadClusterLengths,
+            Sequence<0, 1>,
+            Sequence<0, 1>,
+            I1,
+            ScalarPerVector,
+            Sequence<true>,
+            Sequence<true>>{
+            in_grid_desc,
+            make_tuple(make_multi_index(m_block_data_idx_on_grid, k_block_data_idx_on_grid)),
+            out_grid_desc,
+            make_tuple(make_multi_index(m_block_data_idx_on_grid, k_block_data_idx_on_grid)),
+            tensor_operation::element_wise::PassThrough{}};
+
+        copy_global_to_global.Run(
+            tie(in_grid_desc), tie(in_global_buf), tie(out_grid_desc), tie(out_global_buf));
+    }
+
+    // template <typename... TsIn, typename... TsOut>
+    __host__ static constexpr bool CheckValidity(const InputGridDesc& in_grid_desc,
+                                                 const OutputGridDesc& out_grid_desc)
+    {
+        if(in_grid_desc.GetLength(I0) % MPerBlock != 0 ||
+           in_grid_desc.GetLength(I1) % KPerBlock != 0)
+            return false;
+        if(out_grid_desc.GetLength(I0) % MPerBlock != 0 ||
+           out_grid_desc.GetLength(I1) % KPerBlock != 0)
+            return false;
+        return true;
+    }
+};
+
+} // namespace ck

library/include/ck/library/reference_tensor_operation/cpu/reference_image_to_column.hpp

+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
+
+#pragma once
+
+#include <iostream>
+#include <type_traits>
+#include <sstream>
+
+#include "ck/tensor_operation/gpu/device/device_base.hpp"
+#include "ck/tensor_operation/gpu/device/tensor_layout.hpp"
+#include "ck/library/utility/host_tensor.hpp"
+
+namespace ck {
+namespace tensor_operation {
+namespace host {
+
+//
+// @brief      Reference implementation for image to column.
+//
+// @paragraph
+//             Tensor descriptor in NCHW dimensional order
+//
+// @tparam     InDataType               Input tensor data type.
+// @tparam     OutDataType              Output tensor data type.
+// @tparam     NDimSpatial  Number of spatial dimensions.
+//
+// input descriptor in [N, C, Di, Hi, Wi] order
+// output descriptor in [N * Do * Ho * Wo, C * Z * Y * X] order
+// phyiscal layout is [N, Di, Hi, Wi, C]
+template <ck::index_t NDimSpatial,
+          typename InputLayout,
+          typename InDataType,
+          typename OutDataType,
+          typename std::enable_if<NDimSpatial >= 1 && NDimSpatial <= 3, bool>::type = false>
+struct ReferenceImageToColumn : public device::BaseOperator
+{
+    // Argument
+    struct Argument : public device::BaseArgument
+    {
+        public:
+        Argument(const Tensor<InDataType>& input,
+                 Tensor<OutDataType>& output,
+                 std::vector<ck::index_t> filter_spatial_lengths,
+                 std::vector<ck::index_t> conv_filter_strides,
+                 std::vector<ck::index_t> conv_filter_dilations,
+                 std::vector<ck::index_t> input_left_pads,
+                 std::vector<ck::index_t> input_right_pads)
+            : input_{input},
+              output_{output},
+              conv_strides_{conv_filter_strides},
+              conv_dilations_{conv_filter_dilations},
+              in_left_pads_{input_left_pads},
+              in_right_pads_{input_right_pads},
+              filter_spatial_lengths_{filter_spatial_lengths}
+        {
+            initOutputSpatialLengths();
+        }
+
+        const Tensor<InDataType>& input_;
+        Tensor<OutDataType>& output_;
+
+        std::vector<index_t> conv_strides_;
+        std::vector<index_t> conv_dilations_;
+        std::vector<index_t> in_left_pads_;
+        std::vector<index_t> in_right_pads_;
+
+        std::vector<index_t> filter_spatial_lengths_;
+        std::vector<index_t> output_spatial_lengths_;
+
+        private:
+        void initOutputSpatialLengths()
+        {
+            constexpr auto input_offset_to_spatial = 3;
+
+            for(ck::index_t i = 0; i < NDimSpatial; ++i)
+            {
+                // XEff = (X - 1) * conv_dilation_w + 1;
+                // Wo = (Wi + in_left_pad_w + in_right_pad_w - XEff) / conv_stride_w + 1;
+                const ck::index_t x_eff = (filter_spatial_lengths_[i] - 1) * conv_dilations_[i] + 1;
+
+                output_spatial_lengths_.push_back(
+                    (input_.GetLengths()[i + input_offset_to_spatial] + in_left_pads_[i] +
+                     in_right_pads_[i] - x_eff) /
+                        conv_strides_[i] +
+                    1);
+            }
+        }
+    };
+
+    struct Invoker : public device::BaseInvoker
+    {
+        using Argument = ReferenceImageToColumn::Argument;
+
+        float Run(const Argument& arg)
+        {
+            if(!(arg.input_.GetNumOfDimension() == NDimSpatial + 3 &&
+                 arg.output_.GetNumOfDimension() == 2))
+            {
+                throw std::runtime_error("wrong! inconsistent dimension");
+            }
+
+            const index_t N = arg.input_.GetLengths()[1];
+            const index_t C = arg.input_.GetLengths()[2];
+
+            if constexpr(NDimSpatial == 1)
+            {
+                const index_t Wo = arg.output_spatial_lengths_[0];
+                auto func        = [&](auto n, auto wo) {
+                    index_t row    = n * Wo + wo;
+                    index_t column = 0;
+
+                    for(index_t x = 0; x < arg.filter_spatial_lengths_[0]; ++x)
+                    {
+                        auto wi = static_cast<ck::long_index_t>(wo * arg.conv_strides_[0]) +
+                                  static_cast<ck::long_index_t>(x * arg.conv_dilations_[0]) -
+                                  static_cast<ck::long_index_t>(arg.in_left_pads_[0]);
+
+                        if(wi >= 0 &&
+                           ck::type_convert<std::size_t>(wi) < arg.input_.GetLengths()[3])
+                        {
+                            for(index_t c = 0; c < C; ++c)
+                            {
+                                column++;
+                                InDataType v_in          = arg.input_(0, n, c, wi);
+                                arg.output_(row, column) = ck::type_convert<OutDataType>(v_in);
+                            }
+                        }
+                    }
+                };
+
+                make_ParallelTensorFunctor(func, N, Wo)(std::thread::hardware_concurrency());
+
+                return 0;
+            }
+            else if constexpr(NDimSpatial == 2)
+            {
+                const index_t Ho = arg.output_spatial_lengths_[0];
+                const index_t Wo = arg.output_spatial_lengths_[1];
+
+                auto func = [&](auto n, auto ho, auto wo) {
+                    index_t row    = n * Ho * Wo + ho * Wo + wo;
+                    index_t column = 0;
+
+                    for(index_t y = 0; y < arg.filter_spatial_lengths_[0]; ++y)
+                    {
+                        auto hi = static_cast<ck::long_index_t>(ho * arg.conv_strides_[0]) +
+                                  static_cast<ck::long_index_t>(y * arg.conv_dilations_[0]) -
+                                  static_cast<ck::long_index_t>(arg.in_left_pads_[0]);
+
+                        for(index_t x = 0; x < arg.filter_spatial_lengths_[1]; ++x)
+                        {
+                            auto wi = static_cast<ck::long_index_t>(wo * arg.conv_strides_[1]) +
+                                      static_cast<ck::long_index_t>(x * arg.conv_dilations_[1]) -
+                                      static_cast<ck::long_index_t>(arg.in_left_pads_[1]);
+
+                            if(hi >= 0 &&
+                               ck::type_convert<std::size_t>(hi) < arg.input_.GetLengths()[3] &&
+                               wi >= 0 &&
+                               ck::type_convert<std::size_t>(wi) < arg.input_.GetLengths()[4])
+                            {
+                                for(index_t c = 0; c < C; ++c)
+                                {
+
+                                    InDataType v_in          = arg.input_(0, n, c, hi, wi);
+                                    arg.output_(row, column) = ck::type_convert<OutDataType>(v_in);
+                                    column++;
+                                }
+                            }
+                        }
+                    }
+                };
+
+                make_ParallelTensorFunctor(func, N, Ho, Wo)(std::thread::hardware_concurrency());
+
+                return 0;
+            }
+            else if constexpr(NDimSpatial == 3)
+            {
+                const index_t Do = arg.output_spatial_lengths_[0];
+                const index_t Ho = arg.output_spatial_lengths_[1];
+                const index_t Wo = arg.output_spatial_lengths_[2];
+
+                auto func = [&](auto n, auto d_o, auto ho, auto wo) {
+                    index_t row    = n * Do * Ho * Wo + d_o * Ho * Wo + ho * Wo + wo;
+                    index_t column = 0;
+
+                    for(index_t z = 0; z < arg.filter_spatial_lengths_[0]; ++z)
+                    {
+                        auto di = static_cast<ck::long_index_t>(d_o * arg.conv_strides_[0]) +
+                                  static_cast<ck::long_index_t>(z * arg.conv_dilations_[0]) -
+                                  static_cast<ck::long_index_t>(arg.in_left_pads_[0]);
+                        for(index_t y = 0; y < arg.filter_spatial_lengths_[1]; ++y)
+                        {
+                            auto hi = static_cast<ck::long_index_t>(ho * arg.conv_strides_[1]) +
+                                      static_cast<ck::long_index_t>(y * arg.conv_dilations_[1]) -
+                                      static_cast<ck::long_index_t>(arg.in_left_pads_[1]);
+                            for(index_t x = 0; x < arg.filter_spatial_lengths_[2]; ++x)
+                            {
+                                auto wi =
+                                    static_cast<ck::long_index_t>(wo * arg.conv_strides_[2]) +
+                                    static_cast<ck::long_index_t>(x * arg.conv_dilations_[2]) -
+                                    static_cast<ck::long_index_t>(arg.in_left_pads_[2]);
+                                if(di >= 0 &&
+                                   ck::type_convert<std::size_t>(di) < arg.input_.GetLengths()[3] &&
+                                   hi >= 0 &&
+                                   ck::type_convert<std::size_t>(hi) < arg.input_.GetLengths()[4] &&
+                                   wi >= 0 &&
+                                   ck::type_convert<std::size_t>(wi) < arg.input_.GetLengths()[5])
+                                {
+                                    for(index_t c = 0; c < C; ++c)
+                                    {
+                                        InDataType v_in = arg.input_(0, n, c, di, hi, wi);
+                                        arg.output_(row, column) =
+                                            ck::type_convert<OutDataType>(v_in);
+                                        column++;
+                                    }
+                                }
+                            }
+                        }
+                    }
+                };
+
+                make_ParallelTensorFunctor(func, N, Do, Ho, Wo)(
+                    std::thread::hardware_concurrency());
+
+                return 0;
+            }
+        }
+
+        float Run(const device::BaseArgument* p_arg,
+                  const StreamConfig& /*stream_config*/ = StreamConfig{}) override
+        {
+            return Run(*dynamic_cast<const Argument*>(p_arg));
+        }
+    };
+
+    static constexpr bool IsValidCompilationParameter()
+    {
+        using namespace tensor_layout::convolution;
+
+        if(!(std::is_same_v<InputLayout, GNWC> || std::is_same_v<InputLayout, GNHWC> ||
+             std::is_same_v<InputLayout, GNDHWC>))
+        {
+            return false;
+        }
+        if(!(NDimSpatial >= 1 && NDimSpatial <= 3))
+        {
+            return false;
+        }
+        return true;
+    }
+
+    bool IsSupportedArgument(const Argument& arg)
+    {
+        const ck::index_t G = arg.input_.GetLengths()[0];
+        const ck::index_t N = arg.input_.GetLengths()[1];
+        const ck::index_t C = arg.input_.GetLengths()[2];
+
+        const index_t NDoHoWo =
+            N * ck::accumulate_n<index_t>(
+                    arg.output_spatial_lengths_.begin(), NDimSpatial, 1, std::multiplies<>());
+        const index_t CZYX =
+            C * ck::accumulate_n<index_t>(
+                    arg.filter_spatial_lengths_.begin(), NDimSpatial, 1, std::multiplies<>());
+
+        if(!(arg.output_.GetLengths()[0] == static_cast<std::size_t>(NDoHoWo) &&
+             arg.output_.GetLengths()[1] == static_cast<std::size_t>(CZYX)))
+        {
+            return false;
+        }
+
+        if(G != 1)
+        {
+            return false;
+        }
+        return true;
+    }
+
+    bool IsSupportedArgument(const device::BaseArgument* p_arg) override
+    {
+        return IsSupportedArgument(*dynamic_cast<const Argument*>(p_arg));
+    }
+
+    static auto MakeArgument(const Tensor<InDataType>& input,
+                             Tensor<OutDataType>& output,
+                             std::vector<ck::index_t> filter_spatial_lengths,
+                             std::vector<ck::index_t> conv_filter_strides,
+                             std::vector<ck::index_t> conv_filter_dilations,
+                             std::vector<ck::index_t> input_left_pads,
+                             std::vector<ck::index_t> input_right_pads)
+    {
+        return Argument{input,
+                        output,
+                        filter_spatial_lengths,
+                        conv_filter_strides,
+                        conv_filter_dilations,
+                        input_left_pads,
+                        input_right_pads};
+    }
+
+    static auto MakeInvoker() { return Invoker{}; }
+
+    virtual std::unique_ptr<device::BaseInvoker> MakeInvokerPointer()
+    {
+        return std::make_unique<Invoker>(Invoker{});
+    }
+
+    std::string GetTypeString() const override
+    {
+        auto str = std::stringstream();
+
+        // clang-format off
+        str << "ReferenceImageToColumn"
+            << std::endl;
+        // clang-format on
+
+        return str.str();
+    }
+};
+
+} // namespace host
+} // namespace tensor_operation
+} // namespace ck

library/include/ck/library/tensor_operation_instance/gpu/image_to_column.hpp

+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
+
+#pragma once
+
+#include <vector>
+#include <memory>
+#include "ck/ck.hpp"
+#include "ck/tensor_operation/gpu/device/device_image_to_column.hpp"
+#include "ck/tensor_operation/gpu/device/tensor_layout.hpp"
+
+#include "ck/library/tensor_operation_instance/device_operation_instance_factory.hpp"
+
+namespace ck {
+namespace tensor_operation {
+namespace device {
+namespace instance {
+
+// nhwc, 1d
+void add_device_image_to_column_nhwc_1d_bf16_instances(
+    std::vector<std::unique_ptr<DeviceImageToColumn<1, GNWC, BF16, BF16>>>& instances);
+
+void add_device_image_to_column_nhwc_1d_f16_instances(
+    std::vector<std::unique_ptr<DeviceImageToColumn<1, GNWC, F16, F16>>>& instances);
+
+void add_device_image_to_column_nhwc_1d_f32_instances(
+    std::vector<std::unique_ptr<DeviceImageToColumn<1, GNWC, F32, F32>>>& instances);
+
+void add_device_image_to_column_nhwc_1d_i8_instances(
+    std::vector<std::unique_ptr<DeviceImageToColumn<1, GNWC, int8_t, int8_t>>>& instances);
+// nhwc, 2d
+void add_device_image_to_column_nhwc_2d_bf16_instances(
+    std::vector<std::unique_ptr<DeviceImageToColumn<2, GNHWC, BF16, BF16>>>& instances);
+
+void add_device_image_to_column_nhwc_2d_f16_instances(
+    std::vector<std::unique_ptr<DeviceImageToColumn<2, GNHWC, F16, F16>>>& instances);
+
+void add_device_image_to_column_nhwc_2d_f32_instances(
+    std::vector<std::unique_ptr<DeviceImageToColumn<2, GNHWC, F32, F32>>>& instances);
+
+void add_device_image_to_column_nhwc_2d_i8_instances(
+    std::vector<std::unique_ptr<DeviceImageToColumn<2, GNHWC, int8_t, int8_t>>>& instances);
+// nhwc, 3d
+void add_device_image_to_column_nhwc_3d_bf16_instances(
+    std::vector<std::unique_ptr<DeviceImageToColumn<3, GNDHWC, BF16, BF16>>>& instances);
+
+void add_device_image_to_column_nhwc_3d_f16_instances(
+    std::vector<std::unique_ptr<DeviceImageToColumn<3, GNDHWC, F16, F16>>>& instances);
+
+void add_device_image_to_column_nhwc_3d_f32_instances(
+    std::vector<std::unique_ptr<DeviceImageToColumn<3, GNDHWC, F32, F32>>>& instances);
+
+void add_device_image_to_column_nhwc_3d_i8_instances(
+    std::vector<std::unique_ptr<DeviceImageToColumn<3, GNDHWC, int8_t, int8_t>>>& instances);
+
+template <ck::index_t NumDimSpatial, typename InLayout, typename InDataType, typename OutDataType>
+struct DeviceOperationInstanceFactory<
+    ck::tensor_operation::device::
+        DeviceImageToColumn<NumDimSpatial, InLayout, InDataType, OutDataType>>
+{
+    using DeviceOp = DeviceImageToColumn<NumDimSpatial, InLayout, InDataType, OutDataType>;
+
+    static auto GetInstances()
+    {
+        std::vector<std::unique_ptr<DeviceOp>> op_ptrs;
+
+        if constexpr(NumDimSpatial == 1 && is_same_v<InLayout, GNWC>)
+        {
+            if constexpr(is_same_v<InDataType, float> && is_same_v<OutDataType, float>)
+            {
+                add_device_image_to_column_nhwc_1d_f32_instances(op_ptrs);
+            }
+            else if constexpr(is_same_v<InDataType, half_t> && is_same_v<OutDataType, half_t>)
+            {
+                add_device_image_to_column_nhwc_1d_f16_instances(op_ptrs);
+            }
+            else if constexpr(is_same_v<InDataType, ck::bhalf_t> &&
+                              is_same_v<OutDataType, ck::bhalf_t>)
+            {
+                add_device_image_to_column_nhwc_1d_bf16_instances(op_ptrs);
+            }
+            else if constexpr(is_same_v<InDataType, int8_t> && is_same_v<OutDataType, int8_t>)
+            {
+                add_device_image_to_column_nhwc_1d_i8_instances(op_ptrs);
+            }
+        }
+        else if constexpr(NumDimSpatial == 2 && is_same_v<InLayout, GNHWC>)
+        {
+            if constexpr(is_same_v<InDataType, float> && is_same_v<OutDataType, float>)
+            {
+                add_device_image_to_column_nhwc_2d_f32_instances(op_ptrs);
+            }
+            else if constexpr(is_same_v<InDataType, half_t> && is_same_v<OutDataType, half_t>)
+            {
+                add_device_image_to_column_nhwc_2d_f16_instances(op_ptrs);
+            }
+            else if constexpr(is_same_v<InDataType, ck::bhalf_t> &&
+                              is_same_v<OutDataType, ck::bhalf_t>)
+            {
+                add_device_image_to_column_nhwc_2d_bf16_instances(op_ptrs);
+            }
+            else if constexpr(is_same_v<InDataType, int8_t> && is_same_v<OutDataType, int8_t>)
+            {
+                add_device_image_to_column_nhwc_2d_i8_instances(op_ptrs);
+            }
+        }
+        else if constexpr(NumDimSpatial == 3 && is_same_v<InLayout, GNDHWC>)
+        {
+            if constexpr(is_same_v<InDataType, float> && is_same_v<OutDataType, float>)
+            {
+                add_device_image_to_column_nhwc_3d_f32_instances(op_ptrs);
+            }
+            else if constexpr(is_same_v<InDataType, half_t> && is_same_v<OutDataType, half_t>)
+            {
+                add_device_image_to_column_nhwc_3d_f16_instances(op_ptrs);
+            }
+            else if constexpr(is_same_v<InDataType, ck::bhalf_t> &&
+                              is_same_v<OutDataType, ck::bhalf_t>)
+            {
+                add_device_image_to_column_nhwc_3d_bf16_instances(op_ptrs);
+            }
+            else if constexpr(is_same_v<InDataType, int8_t> && is_same_v<OutDataType, int8_t>)
+            {
+                add_device_image_to_column_nhwc_3d_i8_instances(op_ptrs);
+            }
+        }
+
+        return op_ptrs;
+    }
+};
+
+} // namespace instance
+} // namespace device
+} // namespace tensor_operation
+} // namespace ck

library/include/ck/library/tensor_operation_instance/gpu/image_to_column/device_image_to_column_instance.hpp

+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
+
+#include "ck/ck.hpp"
+#include "ck/tensor_operation/gpu/device/tensor_layout.hpp"
+#include "ck/tensor_operation/gpu/device/impl/device_image_to_column_impl.hpp"
+
+#include "ck/library/tensor_operation_instance/add_device_operation_instance.hpp"
+
+namespace ck {
+namespace tensor_operation {
+namespace device {
+namespace instance {
+
+using namespace ck::tensor_layout::convolution;
+
+using BF16 = ck::bhalf_t;
+using F16  = ck::half_t;
+using F32  = float;
+
+template <ck::index_t... Is>
+using S = ck::Sequence<Is...>;
+
+template <ck::index_t NDimSpatial, typename InLayout>
+using device_image_to_column_bf16_instances = std::tuple<
+    // clang-format off
+        //#####################|        Num| InLayout| InDataType| OutDataType| Block|  MPer|  KPer|       Slice|    Thread| Scalar|
+        //#####################|        Dim|         |           |            |  Size| Block| Block|     Lengths|   Cluster|    Per|
+        //#####################|    Spatial|         |           |            |      |      |      |            |   Lengths| Vector|
+        //#####################|           |         |           |            |      |      |      |            |          |       |
+        DeviceImageToColumnImpl<NDimSpatial, InLayout,       BF16,        BF16,   256,   128,   128, S<128, 128>, S<16, 16>,      8>
+    // clang-format on
+    >;
+
+template <ck::index_t NDimSpatial, typename InLayout>
+using device_image_to_column_f16_instances = std::tuple<
+    // clang-format off
+        //#####################|        Num| InLayout| InDataType| OutDataType| Block|  MPer|  KPer|       Slice|    Thread| Scalar|
+        //#####################|        Dim|         |           |            |  Size| Block| Block|     Lengths|   Cluster|    Per|
+        //#####################|    Spatial|         |           |            |      |      |      |            |   Lengths| Vector|
+        //#####################|           |         |           |            |      |      |      |            |          |       |
+        DeviceImageToColumnImpl<NDimSpatial, InLayout,        F16,         F16,   256,   128,   128, S<128, 128>, S<16, 16>,      8>
+    // clang-format on
+    >;
+
+template <ck::index_t NDimSpatial, typename InLayout>
+using device_image_to_column_f32_instances = std::tuple<
+    // clang-format off
+        //#####################|        Num| InLayout| InDataType| OutDataType| Block|  MPer|  KPer|       Slice|    Thread| Scalar|
+        //#####################|        Dim|         |           |            |  Size| Block| Block|     Lengths|   Cluster|    Per|
+        //#####################|    Spatial|         |           |            |      |      |      |            |   Lengths| Vector|
+        //#####################|           |         |           |            |      |      |      |            |          |       |
+        DeviceImageToColumnImpl<NDimSpatial, InLayout,        F32,         F32,   256,   128,   128, S<128, 128>, S<16, 16>,      4>
+    // clang-format on
+    >;
+
+template <ck::index_t NDimSpatial, typename InLayout>
+using device_image_to_column_i8_instances = std::tuple<
+    // clang-format off
+        //#####################|        Num| InLayout| InDataType| OutDataType| Block|  MPer|  KPer|       Slice|    Thread| Scalar|
+        //#####################|        Dim|         |           |            |  Size| Block| Block|     Lengths|   Cluster|    Per|
+        //#####################|    Spatial|         |           |            |      |      |      |            |   Lengths| Vector|
+        //#####################|           |         |           |            |      |      |      |            |          |       |
+        DeviceImageToColumnImpl<NDimSpatial, InLayout,     int8_t,      int8_t,   256,   256,   256, S<256, 256>, S<16, 16>,     16>
+    // clang-format on
+    >;
+
+} // namespace instance
+} // namespace device
+} // namespace tensor_operation
+} // namespace ck

library/src/tensor_operation_instance/gpu/image_to_column/device_image_to_column_nhwc_1d_instance.cpp

+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
+
+#include "ck/library/tensor_operation_instance/gpu/image_to_column/device_image_to_column_instance.hpp"
+#include "ck/library/tensor_operation_instance/add_device_operation_instance.hpp"
+
+namespace ck {
+namespace tensor_operation {
+namespace device {
+namespace instance {
+
+void add_device_image_to_column_nhwc_1d_bf16_instances(
+    std::vector<std::unique_ptr<DeviceImageToColumn<1, GNWC, BF16, BF16>>>& instances)
+{
+    add_device_operation_instances(instances, device_image_to_column_bf16_instances<1, GNWC>{});
+}
+
+void add_device_image_to_column_nhwc_1d_f16_instances(
+    std::vector<std::unique_ptr<DeviceImageToColumn<1, GNWC, F16, F16>>>& instances)
+{
+    add_device_operation_instances(instances, device_image_to_column_f16_instances<1, GNWC>{});
+}
+
+void add_device_image_to_column_nhwc_1d_f32_instances(
+    std::vector<std::unique_ptr<DeviceImageToColumn<1, GNWC, F32, F32>>>& instances)
+{
+    add_device_operation_instances(instances, device_image_to_column_f32_instances<1, GNWC>{});
+}
+
+void add_device_image_to_column_nhwc_1d_i8_instances(
+    std::vector<std::unique_ptr<DeviceImageToColumn<1, GNWC, int8_t, int8_t>>>& instances)
+{
+    add_device_operation_instances(instances, device_image_to_column_i8_instances<1, GNWC>{});
+}
+
+} // namespace instance
+} // namespace device
+} // namespace tensor_operation
+} // namespace ck

library/src/tensor_operation_instance/gpu/image_to_column/device_image_to_column_nhwc_2d_instance.cpp

+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
+
+#include "ck/library/tensor_operation_instance/gpu/image_to_column/device_image_to_column_instance.hpp"
+#include "ck/library/tensor_operation_instance/add_device_operation_instance.hpp"
+
+namespace ck {
+namespace tensor_operation {
+namespace device {
+namespace instance {
+
+void add_device_image_to_column_nhwc_2d_bf16_instances(
+    std::vector<std::unique_ptr<DeviceImageToColumn<2, GNHWC, BF16, BF16>>>& instances)
+{
+    add_device_operation_instances(instances, device_image_to_column_bf16_instances<2, GNHWC>{});
+}
+
+void add_device_image_to_column_nhwc_2d_f16_instances(
+    std::vector<std::unique_ptr<DeviceImageToColumn<2, GNHWC, F16, F16>>>& instances)
+{
+    add_device_operation_instances(instances, device_image_to_column_f16_instances<2, GNHWC>{});
+}
+
+void add_device_image_to_column_nhwc_2d_f32_instances(
+    std::vector<std::unique_ptr<DeviceImageToColumn<2, GNHWC, F32, F32>>>& instances)
+{
+    add_device_operation_instances(instances, device_image_to_column_f32_instances<2, GNHWC>{});
+}
+
+void add_device_image_to_column_nhwc_2d_i8_instances(
+    std::vector<std::unique_ptr<DeviceImageToColumn<2, GNHWC, int8_t, int8_t>>>& instances)
+{
+    add_device_operation_instances(instances, device_image_to_column_i8_instances<2, GNHWC>{});
+}
+
+} // namespace instance
+} // namespace device
+} // namespace tensor_operation
+} // namespace ck

library/src/tensor_operation_instance/gpu/image_to_column/device_image_to_column_nhwc_3d_instance.cpp

+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
+
+#include "ck/library/tensor_operation_instance/gpu/image_to_column/device_image_to_column_instance.hpp"
+#include "ck/library/tensor_operation_instance/add_device_operation_instance.hpp"
+
+namespace ck {
+namespace tensor_operation {
+namespace device {
+namespace instance {
+
+void add_device_image_to_column_nhwc_3d_bf16_instances(
+    std::vector<std::unique_ptr<DeviceImageToColumn<3, GNDHWC, BF16, BF16>>>& instances)
+{
+    add_device_operation_instances(instances, device_image_to_column_bf16_instances<3, GNDHWC>{});
+}
+
+void add_device_image_to_column_nhwc_3d_f16_instances(
+    std::vector<std::unique_ptr<DeviceImageToColumn<3, GNDHWC, F16, F16>>>& instances)
+{
+    add_device_operation_instances(instances, device_image_to_column_f16_instances<3, GNDHWC>{});
+}
+
+void add_device_image_to_column_nhwc_3d_f32_instances(
+    std::vector<std::unique_ptr<DeviceImageToColumn<3, GNDHWC, F32, F32>>>& instances)
+{
+    add_device_operation_instances(instances, device_image_to_column_f32_instances<3, GNDHWC>{});
+}
+
+void add_device_image_to_column_nhwc_3d_i8_instances(
+    std::vector<std::unique_ptr<DeviceImageToColumn<3, GNDHWC, int8_t, int8_t>>>& instances)
+{
+    add_device_operation_instances(instances, device_image_to_column_i8_instances<3, GNDHWC>{});
+}
+
+} // namespace instance
+} // namespace device
+} // namespace tensor_operation
+} // namespace ck

profiler/src/profile_image_to_column.cpp

+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
+
+#include <iostream>
+#include <numeric>
+#include <initializer_list>
+#include <cstdlib>
+
+#include "profiler/profile_image_to_column_impl.hpp"
+#include "profiler_operation_registry.hpp"
+
+namespace {
+
+enum struct ConvLayout
+{
+    NHWC, // 0
+};
+
+enum struct DataType
+{
+    F32_F32,   // 0
+    F16_F16,   // 1
+    BF16_BF16, // 2
+    INT8_INT8, // 3
+};
+
+#define OP_NAME "image_to_column"
+#define OP_DESC "Image To Column"
+
+static void print_helper_msg()
+{
+    std::cout
+        // clang-format off
+        << "arg1: tensor operation (" OP_NAME ": " OP_DESC ")\n"
+        << "arg2: data type (0: Input fp32, Weight fp32, Output fp32\n"
+        << "                 1: Input fp16, Weight fp16, Output fp16\n"
+        << "                 2: Input bf16, Weight bf16, Output bf16\n"
+        << "                 3: Input int8, Weight int8, Output int8)\n"
+        << "arg3: tensor layout (0: Input[N, Hi, Wi, C], Output[N * Ho * Wo, Y * X * C])\n"
+        << "arg4: verification (0: no, 1: yes)\n"
+        << "arg5: initialization (0: no init, 1: integer value, 2: decimal value)\n"
+        << "arg6: print tensor value (0: no; 1: yes)\n"
+        << "arg7: time kernel (0: no, 1: yes)\n"
+        << ck::utils::conv::get_conv_param_parser_helper_msg() << std::endl;
+    // clang-format on
+}
+
+} // namespace
+
+int profile_grouped_conv_fwd(int argc, char* argv[])
+{
+    // 8 for control, 1 for num_dim_spatial
+    if(argc < 9)
+    {
+        print_helper_msg();
+        return 1;
+    }
+
+    const auto data_type       = static_cast<DataType>(std::stoi(argv[2]));
+    const auto layout          = static_cast<ConvLayout>(std::stoi(argv[3]));
+    const bool do_verification = std::stoi(argv[4]);
+    const int init_method      = std::stoi(argv[5]);
+    const bool do_log          = std::stoi(argv[6]);
+    const bool time_kernel     = std::stoi(argv[7]);
+    const int num_dim_spatial  = std::stoi(argv[8]);
+
+    // 8 for control, 1 for num_dim_spatial, 4 for G/N/K/C, and 6 * num_dim_spatial
+    if(argc != 8 + 1 + 4 + 6 * num_dim_spatial)
+    {
+        print_helper_msg();
+        return 1;
+    }
+
+    const auto params = ck::utils::conv::parse_conv_param(num_dim_spatial, 9, argv);
+
+    using F32  = float;
+    using F16  = ck::half_t;
+    using BF16 = ck::bhalf_t;
+    using INT8 = int8_t;
+
+    using namespace ck::tensor_layout::convolution;
+
+    constexpr auto I1 = ck::Number<1>{};
+    constexpr auto I2 = ck::Number<2>{};
+    constexpr auto I3 = ck::Number<3>{};
+
+    auto profile = [&](auto num_dim_spatial_tmp, auto in_layout, auto in_type, auto out_type) {
+        constexpr ck::index_t NDimSpatial = num_dim_spatial_tmp.value;
+
+        using InLayout = decltype(in_layout);
+
+        using InDataType  = decltype(in_type);
+        using OutDataType = decltype(out_type);
+
+        bool pass = ck::profiler::
+            profile_image_to_column_impl<NDimSpatial, InLayout, InDataType, OutDataType>(
+                do_verification, init_method, do_log, time_kernel, params);
+
+        return pass ? 0 : 1;
+    };
+
+    // NHWC
+    if(layout == ConvLayout::NHWC)
+    {
+        if(num_dim_spatial == 1)
+        {
+            if(data_type == DataType::F32_F32)
+            {
+                return profile(I1, GNWC{}, F32{}, F32{});
+            }
+            else if(data_type == DataType::F16_F16)
+            {
+                return profile(I1, GNWC{}, F16{}, F16{});
+            }
+            else if(data_type == DataType::BF16_BF16)
+            {
+                return profile(I1, GNWC{}, BF16{}, BF16{});
+            }
+            else if(data_type == DataType::INT8_INT8)
+            {
+                return profile(I1, GNWC{}, INT8{}, INT8{});
+            }
+        }
+        else if(num_dim_spatial == 2)
+        {
+            if(data_type == DataType::F32_F32)
+            {
+                return profile(I2, GNHWC{}, F32{}, F32{});
+            }
+            else if(data_type == DataType::F16_F16)
+            {
+                return profile(I2, GNHWC{}, F16{}, F16{});
+            }
+            else if(data_type == DataType::BF16_BF16)
+            {
+                return profile(I2, GNHWC{}, BF16{}, BF16{});
+            }
+            else if(data_type == DataType::INT8_INT8)
+            {
+                return profile(I2, GNHWC{}, INT8{}, INT8{});
+            }
+        }
+        else if(num_dim_spatial == 3)
+        {
+            if(data_type == DataType::F32_F32)
+            {
+                return profile(I3, GNDHWC{}, F32{}, F32{});
+            }
+            else if(data_type == DataType::F16_F16)
+            {
+                return profile(I3, GNDHWC{}, F16{}, F16{});
+            }
+            else if(data_type == DataType::BF16_BF16)
+            {
+                return profile(I3, GNDHWC{}, BF16{}, BF16{});
+            }
+            else if(data_type == DataType::INT8_INT8)
+            {
+                return profile(I3, GNDHWC{}, INT8{}, INT8{});
+            }
+        }
+    }
+
+    std::cout << "this data_type & layout is not implemented" << std::endl;
+
+    return 1;
+}
+
+REGISTER_PROFILER_OPERATION(OP_NAME, OP_DESC, profile_grouped_conv_fwd);