diff --git a/example/ck_tile/37_grouped_convolution/CMakeLists.txt b/example/ck_tile/19_grouped_convolution/CMakeLists.txt similarity index 100% rename from example/ck_tile/37_grouped_convolution/CMakeLists.txt rename to example/ck_tile/19_grouped_convolution/CMakeLists.txt diff --git a/example/ck_tile/37_grouped_convolution/grouped_convolution_forward.cpp b/example/ck_tile/19_grouped_convolution/grouped_convolution_forward.cpp similarity index 98% rename from example/ck_tile/37_grouped_convolution/grouped_convolution_forward.cpp rename to example/ck_tile/19_grouped_convolution/grouped_convolution_forward.cpp index a2c5da56b3..8cdcbaa112 100644 --- a/example/ck_tile/37_grouped_convolution/grouped_convolution_forward.cpp +++ b/example/ck_tile/19_grouped_convolution/grouped_convolution_forward.cpp @@ -20,12 +20,10 @@ template -float grouped_conv_fwd_calc(const ck_tile::GroupedConvHostArgs& args, - const ck_tile::stream_config& s) +float grouped_conv_fwd(const ck_tile::GroupedConvHostArgs& args, const ck_tile::stream_config& s) { constexpr int kBlockPerCu = 1; - // This part comes from the Codegen constexpr ck_tile::index_t M_Tile = 64; constexpr ck_tile::index_t N_Tile = 64; constexpr ck_tile::index_t K_Tile = 64; diff --git a/example/ck_tile/37_grouped_convolution/grouped_convolution_utils.hpp b/example/ck_tile/19_grouped_convolution/grouped_convolution_utils.hpp similarity index 95% rename from example/ck_tile/37_grouped_convolution/grouped_convolution_utils.hpp rename to example/ck_tile/19_grouped_convolution/grouped_convolution_utils.hpp index f514006894..ddaa36605c 100644 --- a/example/ck_tile/37_grouped_convolution/grouped_convolution_utils.hpp +++ b/example/ck_tile/19_grouped_convolution/grouped_convolution_utils.hpp @@ -23,7 +23,10 @@ ck_tile::index_t fill_spatial_dimensions(std::vector& filter_s constexpr ck_tile::index_t non_sp_dims = 3; const ck_tile::index_t n_dim_sp = arg_parser.get_str("in_layout").size() - non_sp_dims; - assert(n_dim_sp >= 1 && n_dim_sp <= 3 && "Wrong layout!"); + if(!(n_dim_sp >= 1 && n_dim_sp <= 3)) + { + throw std::runtime_error("Wrong layout!\n"); + } if(n_dim_sp == 3) { @@ -101,4 +104,4 @@ auto create_args(int argc, char* argv[]) } // host API -float gemm_calc(const ck_tile::GroupedConvHostArgs& args, const ck_tile::stream_config& s); +float grouped_conv_fwd(const ck_tile::GroupedConvHostArgs& args, const ck_tile::stream_config& s); diff --git a/example/ck_tile/37_grouped_convolution/run_grouped_convolution_example.inc b/example/ck_tile/19_grouped_convolution/run_grouped_convolution_example.inc similarity index 98% rename from example/ck_tile/37_grouped_convolution/run_grouped_convolution_example.inc rename to example/ck_tile/19_grouped_convolution/run_grouped_convolution_example.inc index 35c53146c8..c72b17abd6 100644 --- a/example/ck_tile/37_grouped_convolution/run_grouped_convolution_example.inc +++ b/example/ck_tile/19_grouped_convolution/run_grouped_convolution_example.inc @@ -116,8 +116,8 @@ int run_grouped_conv_fwd_example_with_layouts( if(init_method == 0) { - ck_tile::FillUniformDistribution{-1.f, 1.f}(input); - ck_tile::FillUniformDistribution{-1.f, 1.f}(weight); + ck_tile::FillUniformDistribution{-5.f, 5.f}(input); + ck_tile::FillUniformDistribution{-5.f, 5.f}(weight); } else if(init_method == 1) { diff --git a/example/ck_tile/CMakeLists.txt b/example/ck_tile/CMakeLists.txt index af575dd9c2..69d929fbe9 100644 --- a/example/ck_tile/CMakeLists.txt +++ b/example/ck_tile/CMakeLists.txt @@ -18,6 +18,6 @@ add_subdirectory(15_fused_moe) add_subdirectory(16_batched_gemm) add_subdirectory(17_grouped_gemm) add_subdirectory(18_flatmm) +add_subdirectory(19_grouped_convolution) add_subdirectory(35_batched_transpose) add_subdirectory(36_copy) -add_subdirectory(37_grouped_convolution)