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

@@ -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 {

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