navi test

composable_kernel/include/tensor_operation/gridwise_dynamic_gemm_v2.hpp

@@ -51,7 +51,7 @@ struct GridwiseDynamicGemm_km_kn_mn_v3
 {
     __host__ __device__ static constexpr index_t GetSharedMemoryNumberOfByte()
     {
-        const auto E = 4 * 3 * 3;
+        constexpr auto E = EPerBlock * 3 * 3;
 
         constexpr auto max_lds_align =
             math::lcm(Number<ABlockTransferDstScalarPerVector_K>{}, Number<KPerBlock>{});
@@ -84,7 +84,9 @@ struct GridwiseDynamicGemm_km_kn_mn_v3
         constexpr auto I2 = Number<2>{};
         constexpr auto I3 = Number<3>{};
 
-        const auto E = a_e_k_global_desc.GetLength(I0);
+        constexpr auto E = EPerBlock * 3 * 3;
+
+        // const auto E = a_e_k_global_desc.GetLength(I0);
         const auto K = a_e_k_global_desc.GetLength(I1);
 
         const auto N  = b_e_n_ho_wo_global_desc.GetLength(I1);

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

@@ -11,9 +11,9 @@
 #define CK_DEVICE_BACKEND_AMD 1
 
 // GPU ID
-#define CK_AMD_GPU_GFX906 1
+#define CK_AMD_GPU_GFX906 0
 #define CK_AMD_GPU_GFX908 0
-#define CK_AMD_GPU_GFX1030 0
+#define CK_AMD_GPU_GFX1030 1
 
 // HIP version
 #ifndef CK_HIP_VERSION_FLAT

driver/include/device_dynamic_convolution_forward_implicit_gemm_v5r1_nchw_kcyx_nkhw.hpp

@@ -113,22 +113,22 @@ void device_dynamic_convolution_forward_implicit_gemm_v5r1_nchw_kcyx_nkhw(
     constexpr index_t KPerBlock  = 16;
     constexpr index_t HoPerBlock = 8;
     constexpr index_t WoPerBlock = 32;
-    constexpr index_t EPerBlock  = 1;
+    constexpr index_t EPerBlock  = 4;
 
     constexpr index_t KPerThread  = 16;
     constexpr index_t HoPerThread = 2;
     constexpr index_t WoPerThread = 2;
-    constexpr index_t EPerThread  = 1;
+    constexpr index_t EPerThread  = EPerBlock;
 
     using ABlockTransferThreadSliceLengths_E_K   = Sequence<9, 1>;
-    using ABlockTransferThreadClusterLengths_E_K = Sequence<4, 16>;
+    using ABlockTransferThreadClusterLengths_E_K = Sequence<EPerBlock, 16>;
 
     constexpr index_t ABlockTransferSrcScalarPerVector_E = 1;
     constexpr index_t ABlockTransferDstScalarPerVector_K = 1;
 
     constexpr index_t BThreadTransferSrcScalarPerVector_W = 1;
 
-    constexpr index_t CThreadTransferDstScalarPerVector_W = 2;
+    constexpr index_t CThreadTransferDstScalarPerVector_W = 1;
 
     constexpr auto conv_driver =
         DriverDynamicConvolutionForwardImplicitGemm_v5r1_nchw_kcyx_nkhw_pad<

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 -gline-tables-only -save-temps=$CWD -ftemplate-backtrace-limit=0"                   \
+-D CMAKE_CXX_FLAGS="-O3 --amdgpu-target=gfx1030 -gline-tables-only -save-temps=$CWD -ftemplate-backtrace-limit=0"                   \
 -D CMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc                                                                                          \
 -D CMAKE_PREFIX_PATH="/opt/rocm"                                                                                                   \
 -D CMAKE_VERBOSE_MAKEFILE:BOOL=ON                                                                                                  \