From 8a4ca884770845bd5ae7953001f77d74cb0c028b Mon Sep 17 00:00:00 2001
From: Sourabh Betigeri <sourabh.betigeri@amd.com>
Date: Mon, 2 Aug 2021 16:11:51 -0700
Subject: [PATCH] SWDEV-297689 - Moving the vendor specific builtin function
 calls to hip_cooperative_helper.h header file

Change-Id: I797bdde8fbb01285dfda0eac0f960ded260beb8e
---
 hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h | 6 ++----
 .../include/hip/amd_detail/hip_cooperative_groups_helper.h | 7 +++++++
 2 files changed, 9 insertions(+), 4 deletions(-)

diff --git a/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h b/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h
index cd1b0e73b7..65cd8508b6 100644
--- a/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h
+++ b/hipamd/include/hip/amd_detail/amd_hip_cooperative_groups.h
@@ -252,8 +252,7 @@ class tiled_group : public thread_group {
   }
 
   __CG_QUALIFIER__ void sync() const {
-    // enforce memory ordering for memory instructions.
-    __builtin_amdgcn_fence(__ATOMIC_ACQ_REL, "agent");
+    internal::tiled_group::sync();
   }
 };
 
@@ -361,8 +360,7 @@ template <unsigned int size> class thread_block_tile_base : public tile_base<siz
 
  public:
   __CG_STATIC_QUALIFIER__ void sync() {
-    // enforce ordering for memory instructions
-    __builtin_amdgcn_fence(__ATOMIC_ACQ_REL, "agent");
+    internal::tiled_group::sync();
   }
 
   template <class T> __CG_QUALIFIER__ T shfl(T var, int srcRank) const {
diff --git a/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h b/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h
index 3c7068dbb6..e10ae3b9d3 100644
--- a/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h
+++ b/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h
@@ -170,6 +170,13 @@ __CG_STATIC_QUALIFIER__ void sync() { __syncthreads(); }
 
 }  // namespace workgroup
 
+namespace tiled_group {
+
+// enforce ordering for memory intructions
+__CG_STATIC_QUALIFIER__ void sync() { __builtin_amdgcn_fence(__ATOMIC_ACQ_REL, "agent"); }
+
+}  // namespace tiled_group
+
 }  // namespace internal
 
 }  // namespace cooperative_groups