clean

composable_kernel/include/driver/driver_dynamic_convolution_forward_implicit_gemm_v5r1_nchw_kcyx_nkhw.hpp

@@ -24,8 +24,6 @@ template <index_t BlockSize,
           typename GemmABlockTransferThreadClusterLengths_GemmK_GemmM,
           index_t GemmABlockTransferSrcScalarPerVector_GemmK,
           index_t GemmABlockTransferDstScalarPerVector_GemmM,
-          typename GemmBBlockTransferThreadSliceLengths_GemmK_GemmN,
-          typename GemmBBlockTransferThreadClusterLengths_GemmK_GemmN,
           index_t GemmBBlockTransferSrcScalarPerVector_GemmN,
           index_t GemmBBlockTransferDstScalarPerVector_GemmN,
           index_t GemmCThreadTransferDstScalarPerVector_GemmN1>
@@ -191,8 +189,6 @@ struct DriverDynamicConvolutionForwardImplicitGemm_v5r1_nchw_kcyx_nkhw_pad
             GemmABlockTransferSrcScalarPerVector_GemmK,
             GemmABlockTransferDstScalarPerVector_GemmM,
             false, // don't move back src coordinate after threadwise copy
-            GemmBBlockTransferThreadSliceLengths_GemmK_GemmN,
-            GemmBBlockTransferThreadClusterLengths_GemmK_GemmN,
             Sequence<3, 2, 1, 0>,
             Sequence<3, 2, 1, 0>,
             3,

composable_kernel/include/tensor_operation/gridwise_dynamic_gemm_v2.hpp

@@ -34,8 +34,6 @@ template <index_t BlockSize,
           index_t ABlockTransferSrcScalarPerVector,
           index_t ABlockTransferDstScalarPerVector_M,
           bool AThreadTransferSrcResetCoordinateAfterRun,
-          typename BBlockTransferThreadSliceLengths_K_N,
-          typename BBlockTransferThreadClusterLengths_K_N,
           typename BBlockTransferThreadClusterArrangeOrder,
           typename BBlockTransferSrcAccessOrder,
           index_t BBlockTransferSrcVectorDim,

driver/include/device_dynamic_convolution_forward_implicit_gemm_v5r1_nchw_kcyx_nkhw.hpp

@@ -71,13 +71,13 @@ void device_dynamic_convolution_forward_implicit_gemm_v5r1_nchw_kcyx_nkhw(InDesc
     constexpr index_t BlockSize = 64;
 
     constexpr index_t KPerBlock   = 16;
-    constexpr index_t HPerBlock   = 16;
-    constexpr index_t WPerBlock   = 16;
+    constexpr index_t HPerBlock   = 8;
+    constexpr index_t WPerBlock   = 8;
     constexpr index_t CYXPerBlock = 4;
 
-    constexpr index_t KPerThread   = KPerBlock;
-    constexpr index_t HPerThread   = 2;
-    constexpr index_t WPerThread   = 2;
+    constexpr index_t KPerThread   = 8;
+    constexpr index_t HPerThread   = 1;
+    constexpr index_t WPerThread   = 1;
     constexpr index_t CYXPerThread = 4;
 
     using GemmABlockTransferThreadSliceLengths_GemmK_GemmM   = Sequence<1, 1>;
@@ -86,9 +86,6 @@ void device_dynamic_convolution_forward_implicit_gemm_v5r1_nchw_kcyx_nkhw(InDesc
     constexpr index_t GemmABlockTransferSrcScalarPerVector_GemmK = 1;
     constexpr index_t GemmABlockTransferDstScalarPerVector_GemmM = 1;
 
-    using GemmBBlockTransferThreadSliceLengths_GemmK_GemmN   = Sequence<4, 1>;
-    using GemmBBlockTransferThreadClusterLengths_GemmK_GemmN = Sequence<1, 64>;
-
     constexpr index_t GemmBBlockTransferSrcScalarPerVector_GemmN = 1;
     constexpr index_t GemmBBlockTransferDstScalarPerVector_GemmN = 1;
 
@@ -111,8 +108,6 @@ void device_dynamic_convolution_forward_implicit_gemm_v5r1_nchw_kcyx_nkhw(InDesc
             GemmABlockTransferThreadClusterLengths_GemmK_GemmM,
             GemmABlockTransferSrcScalarPerVector_GemmK,
             GemmABlockTransferDstScalarPerVector_GemmM,
-            GemmBBlockTransferThreadSliceLengths_GemmK_GemmN,
-            GemmBBlockTransferThreadClusterLengths_GemmK_GemmN,
             GemmBBlockTransferSrcScalarPerVector_GemmN,
             GemmBBlockTransferDstScalarPerVector_GemmN,
             GemmCThreadTransferDstScalarPerVector_GemmN1>{};

script/cmake-rocm3.7.sh

@@ -10,7 +10,7 @@ cmake
 -D CMAKE_INSTALL_PREFIX=${MY_PROJECT_INSTALL}                                                                                      \
 -D CMAKE_BUILD_TYPE=Release                                                                                                        \
 -D DEVICE_BACKEND="AMD"                                                                                                            \
--D CMAKE_CXX_FLAGS="-O3 --amdgpu-target=gfx906 -save-temps=$CWD -ftemplate-backtrace-limit=0"                   \
+-D CMAKE_CXX_FLAGS="-O3 --amdgpu-target=gfx906 -save-temps=$CWD"                   \
 -D CMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc                                                                                          \
 -D CMAKE_PREFIX_PATH="/opt/rocm"                                                                                                   \
 -D CMAKE_VERBOSE_MAKEFILE:BOOL=ON                                                                                                  \