From 8ba53a00c869ad8ef131359aa6e09b575389261d Mon Sep 17 00:00:00 2001
From: JD <Jehandad.Khan@amd.com>
Date: Fri, 29 Apr 2022 10:36:19 -0500
Subject: [PATCH] Add gfx90a CI stage for tests (#208)

* Add gfx90a CI stage

* upgrade to ROCm 5.1 and fix formatting

[ROCm/composable_kernel commit: 97d8c5045ef102b700878d02ce12b79b8a1e0098]
---
 Dockerfile                                    |  2 +-
 Jenkinsfile                                   | 11 ++++++
 .../gpu/device/device_batched_gemm_xdl.hpp    |  4 +--
 library/src/utility/conv_fwd_util.cpp         | 35 +++++++++----------
 4 files changed, 31 insertions(+), 21 deletions(-)

diff --git a/Dockerfile b/Dockerfile
index fd69a00ee1..c4cf0fac57 100644
--- a/Dockerfile
+++ b/Dockerfile
@@ -1,6 +1,6 @@
 FROM ubuntu:18.04
 
-ARG ROCMVERSION=5.0
+ARG ROCMVERSION=5.1
 ARG OSDB_BKC_VERSION
 
 RUN set -xe
diff --git a/Jenkinsfile b/Jenkinsfile
index 0aeabd690c..824437c970 100644
--- a/Jenkinsfile
+++ b/Jenkinsfile
@@ -235,6 +235,17 @@ pipeline {
                     }
 
                 }
+                stage("Run Tests: gfx90a")
+                {
+                    agent{ label rocmnode("gfx90a")}
+                    environment{
+                        setup_args = """ -D CMAKE_CXX_FLAGS="--offload-arch=gfx90a -O3 " -DBUILD_DEV=On """
+                    }
+                    steps{
+                        buildHipClangJobAndReboot(setup_args:setup_args, config_targets: "check", no_reboot:true, build_type: 'Release')
+                    }
+
+                }
 
             }
         }
diff --git a/include/ck/tensor_operation/gpu/device/device_batched_gemm_xdl.hpp b/include/ck/tensor_operation/gpu/device/device_batched_gemm_xdl.hpp
index 56ec5a7f2c..eda6823424 100644
--- a/include/ck/tensor_operation/gpu/device/device_batched_gemm_xdl.hpp
+++ b/include/ck/tensor_operation/gpu/device/device_batched_gemm_xdl.hpp
@@ -385,8 +385,8 @@ struct DeviceBatchedGemmXdl
               c_grid_desc_m_n_{DeviceBatchedGemmXdl::MakeCGridDescriptor_M_N(M, N, StrideC)},
               c_grid_desc_m0_n0_m1_n1_m2_m3_m4_n2_{},
               compute_ptr_offset_of_batch_{a_grid_desc_k0_m_k1_.GetElementSpaceSize(),
-                                          b_grid_desc_k0_n_k1_.GetElementSpaceSize(),
-                                          c_grid_desc_m_n_.GetElementSpaceSize()},
+                                           b_grid_desc_k0_n_k1_.GetElementSpaceSize(),
+                                           c_grid_desc_m_n_.GetElementSpaceSize()},
               block_2_ctile_map_{},
               M01_{M01},
               N01_{N01},
diff --git a/library/src/utility/conv_fwd_util.cpp b/library/src/utility/conv_fwd_util.cpp
index fde2caa56b..1658450388 100644
--- a/library/src/utility/conv_fwd_util.cpp
+++ b/library/src/utility/conv_fwd_util.cpp
@@ -37,16 +37,16 @@ std::size_t get_flops(ck::index_t N,
 }
 
 ConvParams::ConvParams()
-        : num_dim_spatial(2),
-          N(128),
-          K(256),
-          C(192),
-          filter_spatial_lengths(2, 3),
-          input_spatial_lengths(2, 71),
-          conv_filter_strides(2, 2),
-          conv_filter_dilations(2, 1),
-          input_left_pads(2, 1),
-          input_right_pads(2, 1)
+    : num_dim_spatial(2),
+      N(128),
+      K(256),
+      C(192),
+      filter_spatial_lengths(2, 3),
+      input_spatial_lengths(2, 71),
+      conv_filter_strides(2, 2),
+      conv_filter_dilations(2, 1),
+      input_left_pads(2, 1),
+      input_right_pads(2, 1)
 {
 }
 
@@ -77,9 +77,9 @@ ConvParams::ConvParams(ck::index_t n_dim,
        conv_filter_dilations.size() != num_dim_spatial ||
        input_left_pads.size() != num_dim_spatial || input_right_pads.size() != num_dim_spatial)
     {
-        throw(std::runtime_error(
-            "ConvParams::GetOutputSpatialLengths: "
-            "parameter size is different from number of declared dimensions!"));
+        throw(
+            std::runtime_error("ConvParams::GetOutputSpatialLengths: "
+                               "parameter size is different from number of declared dimensions!"));
     }
 }
 
@@ -91,9 +91,9 @@ std::vector<ck::index_t> ConvParams::GetOutputSpatialLengths() const
        conv_filter_dilations.size() != num_dim_spatial ||
        input_left_pads.size() != num_dim_spatial || input_right_pads.size() != num_dim_spatial)
     {
-        throw(std::runtime_error(
-            "ConvParams::GetOutputSpatialLengths: "
-            "parameter size is different from number of declared dimensions!"));
+        throw(
+            std::runtime_error("ConvParams::GetOutputSpatialLengths: "
+                               "parameter size is different from number of declared dimensions!"));
     }
 
     std::vector<ck::index_t> out_spatial_len(num_dim_spatial, 0);
@@ -101,8 +101,7 @@ std::vector<ck::index_t> ConvParams::GetOutputSpatialLengths() const
     {
         // XEff = (X - 1) * conv_dilation_w + 1;
         // Wo = (Wi + in_left_pad_w + in_right_pad_w - XEff) / conv_stride_w + 1;
-        const ck::index_t idx_eff =
-            (filter_spatial_lengths[i] - 1) * conv_filter_dilations[i] + 1;
+        const ck::index_t idx_eff = (filter_spatial_lengths[i] - 1) * conv_filter_dilations[i] + 1;
         out_spatial_len[i] =
             (input_spatial_lengths[i] + input_left_pads[i] + input_right_pads[i] - idx_eff) /
                 conv_filter_strides[i] +