SWDEV-378321 - Revert "SWDEV-1 - Directly call __builtin_amdgcn_fence"

This reverts commit 1de57c06b3.

Reason for revert: Root cause for regression SWDEV-378321

Change-Id: If3116c5ad1a23db1840e5fc019d2fe997b37d229


[ROCm/clr commit: 5b39c346ad]

projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h

@@ -775,21 +775,21 @@ __device__
 inline
 static void __threadfence()
 {
-    __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "agent");
+  __atomic_work_item_fence(0, __memory_order_seq_cst, __memory_scope_device);
 }
 
 __device__
 inline
 static void __threadfence_block()
 {
-    __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "workgroup");
+  __atomic_work_item_fence(0, __memory_order_seq_cst, __memory_scope_work_group);
 }
 
 __device__
 inline
 static void __threadfence_system()
 {
-    __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "");
+  __atomic_work_item_fence(0, __memory_order_seq_cst, __memory_scope_all_svm_devices);
 }
 
 // abort
@@ -862,11 +862,14 @@ void __assertfail()
 }
 #endif /* defined(_WIN32) || defined(_WIN64) */
 
-__device__ inline static void __work_group_barrier(__cl_mem_fence_flags flags) {
+__device__
+inline
+static void __work_group_barrier(__cl_mem_fence_flags flags, __memory_scope scope)
+{
     if (flags) {
-        __builtin_amdgcn_fence(__ATOMIC_RELEASE, "workgroup");
+        __atomic_work_item_fence(flags, __memory_order_release, scope);
         __builtin_amdgcn_s_barrier();
-        __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "workgroup");
+        __atomic_work_item_fence(flags, __memory_order_acquire, scope);
     } else {
         __builtin_amdgcn_s_barrier();
     }
@@ -876,7 +879,7 @@ __device__
 inline
 static void __barrier(int n)
 {
-  __work_group_barrier((__cl_mem_fence_flags)n);
+  __work_group_barrier((__cl_mem_fence_flags)n, __memory_scope_work_group);
 }
 
 __device__

projects/clr/hipamd/include/hip/amd_detail/device_library_decls.h

@@ -128,4 +128,26 @@ __device__ inline static __local void* __to_local(unsigned x) { return (__local
 #define __CLK_LOCAL_MEM_FENCE    0x01
 typedef unsigned __cl_mem_fence_flags;
 
+typedef enum __memory_scope {
+  __memory_scope_work_item = __OPENCL_MEMORY_SCOPE_WORK_ITEM,
+  __memory_scope_work_group = __OPENCL_MEMORY_SCOPE_WORK_GROUP,
+  __memory_scope_device = __OPENCL_MEMORY_SCOPE_DEVICE,
+  __memory_scope_all_svm_devices = __OPENCL_MEMORY_SCOPE_ALL_SVM_DEVICES,
+  __memory_scope_sub_group = __OPENCL_MEMORY_SCOPE_SUB_GROUP
+} __memory_scope;
+
+// enum values aligned with what clang uses in EmitAtomicExpr()
+typedef enum __memory_order
+{
+  __memory_order_relaxed = __ATOMIC_RELAXED,
+  __memory_order_acquire = __ATOMIC_ACQUIRE,
+  __memory_order_release = __ATOMIC_RELEASE,
+  __memory_order_acq_rel = __ATOMIC_ACQ_REL,
+  __memory_order_seq_cst = __ATOMIC_SEQ_CST
+} __memory_order;
+
+// Linked from hip.amdgcn.bc
+extern "C" __device__ void
+__atomic_work_item_fence(__cl_mem_fence_flags, __memory_order, __memory_scope);
+
 #endif