refactor MultiIndex

composable_kernel/include/kernel_algorithm/dummy_dynamic_transform_v1.hpp

@@ -555,11 +555,7 @@ struct DummyDynamicTransform_v1
 
         MultiIndex<2> idx;
 
-        // initialize idx
-        for(index_t i = 0; i < 2; ++i)
-        {
-            idx(i) = p_wei_global[get_thread_local_1d_id() + i];
-        }
+        static_for<0, 2, 1>{}([&](auto i) { idx(i) = p_wei_global[get_thread_local_1d_id() + i]; });
 
         const index_t niter = p_wei_global[10];
 

composable_kernel/include/kernel_algorithm/dummy_dynamic_transform_v2.hpp

@@ -12,10 +12,10 @@ __host__ __device__ constexpr auto
 map_convolution_into_gemm_v2(const WeiDesc& wei_k_c_y_x_global_desc,
                              const InDesc& in_n_c_hi_wi_global_desc,
                              const OutDesc& out_n_k_ho_wo_global_desc,
-                             const Array<index_t, 2> conv_strides,
-                             const Array<index_t, 2> conv_dilations,
-                             const Array<index_t, 2> in_left_pads,
-                             const Array<index_t, 2> in_right_pads)
+                             const MultiIndex<2> conv_strides,
+                             const MultiIndex<2> conv_dilations,
+                             const MultiIndex<2> in_left_pads,
+                             const MultiIndex<2> in_right_pads)
 {
     constexpr auto i0 = Number<0>{};
     constexpr auto i1 = Number<1>{};
@@ -96,10 +96,10 @@ struct DummyDynamicTransform_v2_1
                           const WeiDesc wei_k_c_y_x_global_desc,
                           const InDesc in_n_c_hi_wi_global_desc,
                           const OutDesc out_n_k_ho_wo_global_desc,
-                          const Array<index_t, 2> conv_strides,
-                          const Array<index_t, 2> conv_dilations,
-                          const Array<index_t, 2> in_left_pads,
-                          const Array<index_t, 2> in_right_pads) const
+                          const MultiIndex<2> conv_strides,
+                          const MultiIndex<2> conv_dilations,
+                          const MultiIndex<2> in_left_pads,
+                          const MultiIndex<2> in_right_pads) const
     {
         const auto transformed_tensor_descs =
             map_convolution_into_gemm_v2(move(wei_k_c_y_x_global_desc),
@@ -124,7 +124,7 @@ struct DummyDynamicTransform_v2_1
             make_dynamic_tensor_coordinate_v2(in_gemmk_gemmn_global_desc, idx);
 
         const auto in_gemmk_gemmn_coord_step = make_dynamic_tensor_coordinate_step_v2(
-            in_gemmk_gemmn_global_desc, MultiIndex<2>{{1, 0}});
+            in_gemmk_gemmn_global_desc, make_multi_index(1, 0));
 
 #pragma unroll
         for(index_t i = 0; i < 10; ++i)
@@ -143,10 +143,10 @@ struct DummyDynamicTransform_v2_1
                           const WeiDesc wei_k_c_y_x_global_desc,
                           const InDesc in_n_c_hi_wi_global_desc,
                           const OutDesc out_n_k_ho_wo_global_desc,
-                          const Array<index_t, 2> conv_strides,
-                          const Array<index_t, 2> conv_dilations,
-                          const Array<index_t, 2> in_left_pads,
-                          const Array<index_t, 2> in_right_pads) const
+                          const MultiIndex<2> conv_strides,
+                          const MultiIndex<2> conv_dilations,
+                          const MultiIndex<2> in_left_pads,
+                          const MultiIndex<2> in_right_pads) const
     {
         const index_t N = in_n_c_hi_wi_global_desc.GetLength(0);
         const index_t C = in_n_c_hi_wi_global_desc.GetLength(1);
@@ -262,10 +262,10 @@ struct DummyDynamicTransform_v2_1
                         const WeiDesc wei_k_c_y_x_global_desc,
                         const InDesc in_n_c_hi_wi_global_desc,
                         const OutDesc out_n_k_ho_wo_global_desc,
-                        const Array<index_t, 2> conv_strides,
-                        const Array<index_t, 2> conv_dilations,
-                        const Array<index_t, 2> in_left_pads,
-                        const Array<index_t, 2> in_right_pads) const
+                        const MultiIndex<2> conv_strides,
+                        const MultiIndex<2> conv_dilations,
+                        const MultiIndex<2> in_left_pads,
+                        const MultiIndex<2> in_right_pads) const
     {
         Run_1(p_wei_global,
               p_in_global,

composable_kernel/include/tensor_description/dynamic_multi_index_transform.hpp

@@ -279,7 +279,8 @@ struct DynamicMerge
         : low_lengths_{low_lengths},
           low_lengths_scan_{reverse_exclusive_scan_on_array(
               low_lengths, math::multiplies<index_t>{}, index_t{1})},
-          up_lengths_{{reduce_on_array(low_lengths, math::multiplies<index_t>(), index_t{1})}}
+          up_lengths_{make_multi_index(
+              reduce_on_array(low_lengths, math::multiplies<index_t>(), index_t{1}))}
     {
         static_assert(LowerIndex::Size() == NDimLow, "wrong!");
     }

composable_kernel/include/tensor_description/dynamic_tensor_descriptor_v2.hpp

@@ -178,10 +178,10 @@ struct DynamicTensorDescriptor_v2
                                                                       index_t element_space_size)
     {
         // zero initialization
-        HiddenIndex hidden_lengths{{0}};
+        HiddenIndex hidden_lengths = make_zero_multi_index<ndim_hidden_>();
 
         // this is the orignal tensor element space size
-        hidden_lengths(0) = element_space_size;
+        hidden_lengths(Number<0>{}) = element_space_size;
 
         // lengths for all other hidden dimensions
         static_for<0, ntransform_, 1>{}([&transforms, &hidden_lengths](auto itran) {

composable_kernel/include/tensor_description/multi_index.hpp

@@ -5,7 +5,7 @@
 
 namespace ck {
 
-#if 1 // debug
+#if 1
 template <index_t N>
 using MultiIndex = Array<index_t, N>;
 

composable_kernel/include/utility/array.hpp

@@ -51,6 +51,13 @@ __host__ __device__ constexpr auto make_array(const X& x, const Xs&... xs)
     return Array<X, sizeof...(Xs) + 1>{{x, static_cast<X>(xs)...}};
 }
 
+// make empty array
+template <typename X>
+__host__ __device__ constexpr auto make_array()
+{
+    return Array<X, 0>{};
+}
+
 template <typename TData, index_t NSize>
 __host__ __device__ constexpr auto push_back(Array<TData, NSize>& a, const TData& x)
 {

composable_kernel/include/utility/statically_indexed_array.hpp

@@ -17,6 +17,8 @@ struct StaticallyIndexedArray<TData, 0> : Tuple<>
 {
     using data_type = TData;
     using base      = Tuple<>;
+
+    __host__ __device__ explicit constexpr StaticallyIndexedArray() : base() {}
 };
 
 template <typename TData>
@@ -387,5 +389,12 @@ __host__ __device__ constexpr auto make_statically_indexed_array(const X& x, con
     return StaticallyIndexedArray<X, sizeof...(Xs) + 1>(x, static_cast<X>(xs)...);
 }
 
+// make empty StaticallyIndexedArray
+template <typename X>
+__host__ __device__ constexpr auto make_statically_indexed_array()
+{
+    return StaticallyIndexedArray<X, 0>();
+}
+
 } // namespace ck
 #endif

driver/include/device_dummy_dynamic_transform_v2.hpp

@@ -101,10 +101,10 @@ void device_dummy_dynamic_transform_v2(InDesc,
                                                  const decltype(wei_kcyx_desc),
                                                  const decltype(in_nchw_desc),
                                                  const decltype(out_nkhw_desc),
-                                                 const Array<index_t, 2>,
-                                                 const Array<index_t, 2>,
-                                                 const Array<index_t, 2>,
-                                                 const Array<index_t, 2>>,
+                                                 const MultiIndex<2>,
+                                                 const MultiIndex<2>,
+                                                 const MultiIndex<2>,
+                                                 const MultiIndex<2>>,
                           dim3(GridSize),
                           dim3(BlockSize),
                           0,
@@ -125,10 +125,10 @@ void device_dummy_dynamic_transform_v2(InDesc,
                                                  float* const,
                                                  float* const,
                                                  const decltype(in_gemmk_gemmn_global_desc),
-                                                 const Array<index_t, 2>,
-                                                 const Array<index_t, 2>,
-                                                 const Array<index_t, 2>,
-                                                 const Array<index_t, 2>>,
+                                                 const MultiIndex<2>,
+                                                 const MultiIndex<2>,
+                                                 const MultiIndex<2>,
+                                                 const MultiIndex<2>>,
                           dim3(GridSize),
                           dim3(BlockSize),
                           0,