rename

composable_kernel/include/kernel_algorithm/gridwise_convolution_backward_data_implicit_gemm_v1_nchw_kcyx_nkhw_lds_double_buffer.hpp → composable_kernel/include/kernel_algorithm/gridwise_convolution_backward_data_implicit_gemm_v4r4_nchw_kcyx_nkhw_lds_double_buffer.hpp

-#ifndef CK_GRIDWISE_CONVOLUTION_BACKWARD_DATA_IMPLICIT_GEMM_V1_NCHW_KCYX_NKHW_LDS_DOUBLE_BUFFER_HPP
-#define CK_GRIDWISE_CONVOLUTION_BACKWARD_DATA_IMPLICIT_GEMM_V1_NCHW_KCYX_NKHW_LDS_DOUBLE_BUFFER_HPP
+#ifndef CK_GRIDWISE_CONVOLUTION_BACKWARD_DATA_IMPLICIT_GEMM_V4R4_NCHW_KCYX_NKHW_LDS_DOUBLE_BUFFER_HPP
+#define CK_GRIDWISE_CONVOLUTION_BACKWARD_DATA_IMPLICIT_GEMM_V4R4_NCHW_KCYX_NKHW_LDS_DOUBLE_BUFFER_HPP
 
 #include "common_header.hpp"
 #include "tensor_descriptor.hpp"
@@ -41,7 +41,7 @@ template <index_t GridSize,
           typename WeiBlockCopyClusterLengths_K_E,
           index_t WeiBlockCopyDataPerAccess_E,
           index_t InThreadCopyDataPerAccess_B>
-struct GridwiseConvolutionBackwardDataImplicitGemm_v1_nchw_kcyx_nkhw_lds_double_buffer
+struct GridwiseConvolutionBackwardDataImplicitGemm_v4r4_nchw_kcyx_nkhw_lds_double_buffer
 {
     __device__ void Run(Float* const __restrict__ p_in_global,
                         const Float* const __restrict__ p_wei_global,

driver/include/device_convolution_bwd_data_implicit_gemm_v1_nchw_kcyx_nkhw.hpp → driver/include/device_convolution_backward_data_implicit_gemm_v4r4_nchw_kcyx_nkhw.hpp

@@ -3,7 +3,7 @@
 #include "device.hpp"
 #include "tensor.hpp"
 #include "gridwise_operation_wrapper.hpp"
-#include "gridwise_convolution_backward_data_implicit_gemm_v1_nchw_kcyx_nkhw_lds_double_buffer.hpp"
+#include "gridwise_convolution_backward_data_implicit_gemm_v4r4_nchw_kcyx_nkhw_lds_double_buffer.hpp"
 
 template <typename T,
           typename InDesc,
@@ -13,7 +13,7 @@ template <typename T,
           typename ConvDilations,
           typename LeftPads,
           typename RightPads>
-void device_convolution_bwd_data_implicit_gemm_v1_nchw_kcyx_nkhw(InDesc in_nchw_desc,
+void device_convolution_backward_data_implicit_gemm_v4r4_nchw_kcyx_nkhw(InDesc in_nchw_desc,
                                                                         Tensor<T>& in_nchw,
                                                                         WeiDesc wei_kcyx_desc,
                                                                         const Tensor<T>& wei_kcyx,
@@ -85,7 +85,7 @@ void device_convolution_bwd_data_implicit_gemm_v1_nchw_kcyx_nkhw(InDesc in_nchw_
     printf("%s: BlockSize %u, GridSize %u \n", __func__, BlockSize, GridSize);
 
     constexpr auto gridwise_conv =
-        GridwiseConvolutionBackwardDataImplicitGemm_v1_nchw_kcyx_nkhw_lds_double_buffer<
+        GridwiseConvolutionBackwardDataImplicitGemm_v4r4_nchw_kcyx_nkhw_lds_double_buffer<
             GridSize,
             BlockSize,
             T,

driver/src/conv_bwd_data_driver.cpp

@@ -13,7 +13,7 @@
 #include "device_tensor.hpp"
 #include "conv_common.hpp"
 #include "host_conv_bwd_data.hpp"
-#include "device_convolution_bwd_data_implicit_gemm_v1_nchw_kcyx_nkhw.hpp"
+#include "device_convolution_backward_data_implicit_gemm_v4r4_nchw_kcyx_nkhw.hpp"
 
 int main(int argc, char* argv[])
 {
@@ -96,7 +96,7 @@ int main(int argc, char* argv[])
 
     using LeftPads  = Sequence<0, 0>;
     using RightPads = Sequence<0, 0>;
-#elif 0
+#elif 1
     // 1x1 filter, 8x8 image
     // cudnn@V100 83%, ck@V100 75%, ck@P100 78%, ck@VII 65%
     constexpr index_t N  = 128;
@@ -344,8 +344,7 @@ int main(int argc, char* argv[])
 #endif
     }
 
-#if 1
-    device_convolution_bwd_data_implicit_gemm_v1_nchw_kcyx_nkhw(in_nchw_desc,
+    device_convolution_backward_data_implicit_gemm_v4r4_nchw_kcyx_nkhw(in_nchw_desc,
                                                                        in_nchw_device,
                                                                        wei_kcyx_desc,
                                                                        wei_kcyx,
@@ -356,7 +355,6 @@ int main(int argc, char* argv[])
                                                                        LeftPads{},
                                                                        RightPads{},
                                                                        nrepeat);
-#endif
 
     if(do_verification)
     {