tweak

composable_kernel/include/tensor_operation/gridwise_dynamic_gemm_v2.hpp

@@ -102,6 +102,8 @@ struct GridwiseDynamicGemm_km_kn_mn_v3
         const index_t k_block_work_id   = get_block_1d_id() / hwo_block_work_num;
         const index_t hwo_block_work_id = get_block_1d_id() - k_block_work_id * hwo_block_work_num;
 
+        const index_t ho_block_work_id = hwo_block_work_id / wo_block_work_num;
+        const index_t wo_block_work_id = hwo_block_work_id - ho_block_work_id * wo_block_work_num;
 #else
         // Hack: this force result into SGPR
         const index_t k_block_work_num   = __builtin_amdgcn_readfirstlane(K / KPerBlock);
@@ -112,10 +114,11 @@ struct GridwiseDynamicGemm_km_kn_mn_v3
         const index_t k_block_work_id =
             __builtin_amdgcn_readfirstlane(get_block_1d_id() / hwo_block_work_num);
         const index_t hwo_block_work_id = get_block_1d_id() - k_block_work_id * hwo_block_work_num;
-#endif
 
-        const index_t ho_block_work_id = hwo_block_work_id / wo_block_work_num;
+        const index_t ho_block_work_id =
+            __builtin_amdgcn_readfirstlane(hwo_block_work_id / wo_block_work_num);
         const index_t wo_block_work_id = hwo_block_work_id - ho_block_work_id * wo_block_work_num;
+#endif
 
         // lds max alignment
         constexpr auto max_lds_align =

composable_kernel/include/utility/config.amd.hpp.in

@@ -85,7 +85,7 @@
 
 // experimental implementation
 #ifndef CK_EXPERIMENTAL_USE_BUFFER_LOAD_OOB_CHECK_OFFSET_TRICK
-#define CK_EXPERIMENTAL_USE_BUFFER_LOAD_OOB_CHECK_OFFSET_TRICK 1
+#define CK_EXPERIMENTAL_USE_BUFFER_LOAD_OOB_CHECK_OFFSET_TRICK 0
 #endif
 
 #ifndef CK_EXPERIMENTAL_USE_BUFFER_STORE_OOB_CHECK_OFFSET_TRICK

driver/include/device_dynamic_convolution_forward_implicit_gemm_v5r1_nchw_kcyx_nkhw.hpp

@@ -76,8 +76,8 @@ void device_dynamic_convolution_forward_implicit_gemm_v5r1_nchw_kcyx_nkhw(InDesc
     constexpr index_t EPerBlock  = 2;
 
     constexpr index_t KPerThread  = 16;
-    constexpr index_t HoPerThread = 4;
-    constexpr index_t WoPerThread = 1;
+    constexpr index_t HoPerThread = 2;
+    constexpr index_t WoPerThread = 2;
     constexpr index_t EPerThread  = 2;
 
     using GemmABlockTransferThreadSliceLengths_GemmK_GemmM   = Sequence<9, 1>;

driver/src/conv_driver.cpp

@@ -48,8 +48,8 @@ int main(int argc, char* argv[])
     using ConvStrides   = Sequence<1, 1>;
     using ConvDilations = Sequence<1, 1>;
 
-    using LeftPads  = Sequence<0, 0>;
-    using RightPads = Sequence<0, 0>;
+    using LeftPads                   = Sequence<0, 0>;
+    using RightPads                  = Sequence<0, 0>;
 #elif 0
     constexpr index_t N  = 1;
     constexpr index_t C  = 16;
@@ -719,7 +719,7 @@ int main(int argc, char* argv[])
                                                                  LeftPads{},
                                                                  RightPads{},
                                                                  nrepeat);
-#elif 1
+#elif 0
     device_dynamic_convolution_forward_implicit_gemm_v4r4_nchw_kcyx_nkhw<in_data_t,
                                                                          in_vector_size,
                                                                          acc_data_t,
@@ -736,7 +736,7 @@ int main(int argc, char* argv[])
          LeftPads{},
          RightPads{},
          nrepeat);
-#elif 1
+#elif 0
     device_dynamic_convolution_forward_implicit_gemm_v4r4_nhwc_kyxc_nhwk<in_data_t,
                                                                          in_vector_size,
                                                                          acc_data_t,