clean up

composable_kernel/include/kernel_algorithm/gridwise_convolution_implicit_gemm_v4r1_nchw_kcyx_nkhw_lds_double_buffer.hpp

@@ -49,8 +49,7 @@ template <index_t GridSize,
           index_t WeiBlockCopyDstDataPerWrite_K>
 struct GridwiseConvolutionImplicitGemm_v4r1_nchw_kcyx_nkhw_lds_double_buffer
 {
-    __device__ void __launch_bounds__(BlockSize, 2)
-        Run(const Float* const __restrict__ p_in_global,
+    __device__ void Run(const Float* const __restrict__ p_in_global,
                         const Float* const __restrict__ p_wei_global,
                         Float* const __restrict__ p_out_global) const
     {
@@ -268,7 +267,7 @@ struct GridwiseConvolutionImplicitGemm_v4r1_nchw_kcyx_nkhw_lds_double_buffer
         // c_thread_mtx definition: this is a mess
         // TODO:: more elegent way of defining c_thread_mtx
         constexpr auto c_k0k2_n1n2_thread_mtx_desc = make_ConstantMatrixDescriptor_packed(
-            Number<GemmMRepeat * GemmMPerThreadSubC>{}, Number<N1 * N2>{});
+            Number<GemmMRepeat * GemmMPerThreadSubC>{}, Number<GemmNRepeat * GemmMPerThreadSubC>{});
 
         const auto blockwise_gemm = BlockwiseGemmBlockABlockBThreadCTransANormalBNormalC_v2<
             BlockSize,

composable_kernel/include/tensor_description/tensor_coordinate.hpp

@@ -234,8 +234,6 @@ struct MergedTensorCoordinate
     {
         static_assert(is_same<typename T::data_type, index_t>{} && T::GetSize() == nDim, "wrong!");
 
-        index_t normal_offset_diff = 0;
-
         static_for<0, nDim, 1>{}([&](auto idim) {
             if(step_sizes[idim] != 0)
             {

composable_kernel/include/tensor_operation/threadwise_generic_tensor_slice_copy.hpp

@@ -198,7 +198,7 @@ struct ThreadwiseGenericTensorSliceCopy_v1r1
                     dst_vector_access_dim,
                     dst_access_id[dst_vector_access_dim] * dst_data_per_access);
 
-                vector_t vector_data;
+                vector_t vector_data{};
 
                 // pack vector from buffer
                 static_for<0, DstDataPerAccess, 1>{}([&](auto i) {
@@ -224,7 +224,7 @@ struct ThreadwiseGenericTensorSliceCopy_v1r1
                 dst_data_begin_id(dst_vector_access_dim) =
                     dst_access_id[dst_vector_access_dim] * dst_data_per_access;
 
-                vector_t vector_data;
+                vector_t vector_data{};
 
                 // pack vector from buffer
                 for(index_t i = 0; i < DstDataPerAccess; ++i)

driver/include/device_convolution_implicit_gemm_v4r1_nchw_kcyx_nkhw.hpp

@@ -91,8 +91,6 @@ void device_convolution_implicit_gemm_v4r1_nchw_kcyx_nkhw(InDesc,
 
     constexpr index_t WeiBlockCopySrcDataPerRead_E  = 4;
     constexpr index_t WeiBlockCopyDstDataPerWrite_K = 1;
-
-    constexpr index_t OutThreadCopyDataPerAccess_W = 1;
 #elif 1
     // each thread hold 64 data
     constexpr index_t BlockSize = 256;
@@ -101,6 +99,8 @@ void device_convolution_implicit_gemm_v4r1_nchw_kcyx_nkhw(InDesc,
     constexpr index_t KPerBlock = 128;
     constexpr index_t EPerBlock = 8;
 
+    constexpr index_t GemmNRepeat = 2;
+
     constexpr index_t GemmMPerThreadSubC = 4;
     constexpr index_t GemmNPerThreadSubC = 4;
     constexpr index_t GemmMLevel0Cluster = 4;
@@ -136,6 +136,8 @@ void device_convolution_implicit_gemm_v4r1_nchw_kcyx_nkhw(InDesc,
     constexpr index_t KPerBlock = 64;
     constexpr index_t EPerBlock = 8;
 
+    constexpr index_t GemmNRepeat = 2;
+
     constexpr index_t GemmMPerThreadSubC = 2;
     constexpr index_t GemmNPerThreadSubC = 4;
     constexpr index_t GemmMLevel0Cluster = 4;