From ae1aed5111d1e8d2e71f44133e8f94726ebb22c5 Mon Sep 17 00:00:00 2001
From: zjing14 <zhangjing14@gmail.com>
Date: Fri, 26 Aug 2022 09:59:50 -0500
Subject: [PATCH] Fixed splitk gemm fp32 (#384)

* add scripts

* fixed splitK_gemm_fp32

* clean

* clean

[ROCm/composable_kernel commit: 9881625b2d90b897f8c88e0940f8fab657293d0d]
---
 .../gpu/device/device_gemm_xdl_splitk.hpp                 | 4 ++--
 script/run_full_performance_tests.sh                      | 8 ++++----
 2 files changed, 6 insertions(+), 6 deletions(-)

diff --git a/include/ck/tensor_operation/gpu/device/device_gemm_xdl_splitk.hpp b/include/ck/tensor_operation/gpu/device/device_gemm_xdl_splitk.hpp
index b5eed11aeb..62832c3a71 100644
--- a/include/ck/tensor_operation/gpu/device/device_gemm_xdl_splitk.hpp
+++ b/include/ck/tensor_operation/gpu/device/device_gemm_xdl_splitk.hpp
@@ -93,9 +93,9 @@ struct DeviceGemmXdlSplitK : public DeviceGemmSplitK<ALayout,
 
         const auto a_grid_desc_m_kpad = transform_tensor_descriptor(
             a_grid_desc_m_k,
-            make_tuple(make_right_pad_transform(K, KPad - K), make_pass_through_transform(M)),
+            make_tuple(make_pass_through_transform(M), make_right_pad_transform(K, KPad - K)),
             make_tuple(Sequence<0>{}, Sequence<1>{}),
-            make_tuple(Sequence<0>{}, Sequence<1>{}));
+            make_tuple(Sequence<1>{}, Sequence<0>{}));
 
         if constexpr(GemmSpec == GemmSpecialization::MNPadding)
         {
diff --git a/script/run_full_performance_tests.sh b/script/run_full_performance_tests.sh
index bd2d48b668..be90d84c78 100755
--- a/script/run_full_performance_tests.sh
+++ b/script/run_full_performance_tests.sh
@@ -127,10 +127,10 @@ print_log_header $reduction_log $env_type $branch $host_name
 export splitK_gemm_log="perf_splitK_gemm_${gpu_arch}.log"
 print_log_header $splitK_gemm_log $env_type $branch $host_name
 
-#../script/profile_splitK_gemm.sh gemm_splitk 0 0 $verify 1 0 1 4 | tee -a $splitK_gemm_log
-#../script/profile_splitK_gemm.sh gemm_splitk 0 1 $verify 1 0 1 4 | tee -a $splitK_gemm_log
-#../script/profile_splitK_gemm.sh gemm_splitk 0 2 $verify 1 0 1 4 | tee -a $splitK_gemm_log
-#../script/profile_splitK_gemm.sh gemm_splitk 0 3 $verify 1 0 1 4 | tee -a $splitK_gemm_log
+../script/profile_splitK_gemm.sh gemm_splitk 0 0 $verify 1 0 1 4 | tee -a $splitK_gemm_log
+../script/profile_splitK_gemm.sh gemm_splitk 0 1 $verify 1 0 1 4 | tee -a $splitK_gemm_log
+../script/profile_splitK_gemm.sh gemm_splitk 0 2 $verify 1 0 1 4 | tee -a $splitK_gemm_log
+../script/profile_splitK_gemm.sh gemm_splitk 0 3 $verify 1 0 1 4 | tee -a $splitK_gemm_log
 
 ../script/profile_splitK_gemm.sh gemm_splitk 1 0 $verify 1 0 1 4 | tee -a $splitK_gemm_log
 ../script/profile_splitK_gemm.sh gemm_splitk 1 1 $verify 1 0 1 4 | tee -a $splitK_gemm_log