debugged bwd data v2r1

composable_kernel/include/kernel_algorithm/gridwise_convolution_backward_data_implicit_gemm_v2r1_nchw_kcyx_nkhw.hpp

@@ -91,9 +91,22 @@ struct GridwiseConvolutionBackwardDataImplicitGemm_v2r1_nchw_kcyx_nkhw
         constexpr index_t Xdot = math::integer_divide_ceil(X, Xtilda);
 
         constexpr index_t Htilda =
-            Ho + math::integer_divide_ceil(ConvDilationH * (Y - (Y % Ytilda) - 1), ConvStrideH);
+            Ho + math::integer_divide_ceil(ConvDilationH * (Y - 1), ConvStrideH);
         constexpr index_t Wtilda =
-            Wo + math::integer_divide_ceil(ConvDilationW * (X - (X % Xtilda) - 1), ConvStrideW);
+            Wo + math::integer_divide_ceil(ConvDilationW * (X - 1), ConvStrideW);
+
+        constexpr index_t HtildaLeft = math::integer_divide_floor(
+            math::max(0, InLeftPads{}[0] - ConvDilationH * (Ytilda - 1)), ConvStrides{}[0]);
+        constexpr index_t WtildaLeft = math::integer_divide_floor(
+            math::max(0, InLeftPads{}[1] - ConvDilationW * (Xtilda - 1)), ConvStrides{}[1]);
+
+        constexpr index_t HtildaRight = math::min(
+            Htilda, math::integer_divide_ceil(InLeftPads{}[0] + Hi - 1, ConvStrides{}[0]) + 1);
+        constexpr index_t WtildaRight = math::min(
+            Wtilda, math::integer_divide_ceil(InLeftPads{}[1] + Wi - 1, ConvStrides{}[1]) + 1);
+
+        constexpr index_t HtildaTrim = HtildaRight - HtildaLeft;
+        constexpr index_t WtildaTrim = WtildaRight - WtildaLeft;
 
         // weight tensor
         constexpr auto wei_k_c_ydot_ytilda_xdot_xtilda_global_desc = transform_tensor_descriptor(
@@ -115,19 +128,6 @@ struct GridwiseConvolutionBackwardDataImplicitGemm_v2r1_nchw_kcyx_nkhw
             make_tuple(Sequence<0, 2, 4>{}, Sequence<1, 3, 5>{}),
             make_tuple(Sequence<0>{}, Sequence<1>{}));
 
-        constexpr index_t HtildaLeft = math::integer_divide_floor(
-            math::max(0, InLeftPads{}[0] - ConvDilationH * (Ytilda - 1)), ConvStrides{}[0]);
-        constexpr index_t WtildaLeft = math::integer_divide_floor(
-            math::max(0, InLeftPads{}[1] - ConvDilationW * (Xtilda - 1)), ConvStrides{}[1]);
-
-        constexpr index_t HtildaRight = math::min(
-            Htilda, math::integer_divide_ceil(InLeftPads{}[0] + Hi - 1, ConvStrides{}[0]) + 1);
-        constexpr index_t WtildaRight = math::min(
-            Wtilda, math::integer_divide_ceil(InLeftPads{}[1] + Wi - 1, ConvStrides{}[1]) + 1);
-
-        constexpr index_t HtildaTrim = HtildaRight - HtildaLeft;
-        constexpr index_t WtildaTrim = WtildaRight - WtildaLeft;
-
         // output tensor
         constexpr auto out_n_k_ydot_htilda_xdot_wtilda_global_desc = transform_tensor_descriptor(
             out_n_k_ho_wo_global_desc,

driver/include/device_convolution_backward_data_implicit_gemm_v2r1_nchw_kcyx_nkhw.hpp

@@ -157,10 +157,8 @@ void device_convolution_backward_data_implicit_gemm_v2r1_nchw_kcyx_nkhw(InDesc i
     constexpr index_t Ydot = math::integer_divide_ceil(Y, Ytilda);
     constexpr index_t Xdot = math::integer_divide_ceil(X, Xtilda);
 
-    constexpr index_t Htilda =
-        Ho + math::integer_divide_ceil(ConvDilationH * (Y - (Y % Ytilda) - 1), ConvStrideH);
-    constexpr index_t Wtilda =
-        Wo + math::integer_divide_ceil(ConvDilationW * (X - (X % Xtilda) - 1), ConvStrideW);
+    constexpr index_t Htilda = Ho + math::integer_divide_ceil(ConvDilationH * (Y - 1), ConvStrideH);
+    constexpr index_t Wtilda = Wo + math::integer_divide_ceil(ConvDilationW * (X - 1), ConvStrideW);
 
     constexpr index_t HtildaLeft = math::integer_divide_floor(
         math::max(0, InLeftPads{}[0] - ConvDilationH * (Ytilda - 1)), ConvStrides{}[0]);