Revert "[OpenMP] Replace AMDGPU fences with generic scoped fences (#119619)"

This reverts commit f4ee5a6.

Author: jsji

Diff for: offload/DeviceRTL/include/Synchronization.h

@@ -26,14 +26,6 @@ enum OrderingTy {
   seq_cst = __ATOMIC_SEQ_CST,
 };
 
-enum ScopeTy {
-  system = __MEMORY_SCOPE_SYSTEM,
-  device_ = __MEMORY_SCOPE_DEVICE,
-  workgroup = __MEMORY_SCOPE_WRKGRP,
-  wavefront = __MEMORY_SCOPE_WVFRNT,
-  single = __MEMORY_SCOPE_SINGLE,
-};
-
 enum MemScopeTy {
   all,    // All threads on all devices
   device, // All threads on the device

Diff for: offload/DeviceRTL/src/Synchronization.cpp

@@ -232,16 +232,50 @@ void namedBarrier() {
   fence::team(atomic::release);
 }
 
+// sema checking of amdgcn_fence is aggressive. Intention is to patch clang
+// so that it is usable within a template environment and so that a runtime
+// value of the memory order is expanded to this switch within clang/llvm.
 void fenceTeam(atomic::OrderingTy Ordering) {
-  return __scoped_atomic_thread_fence(Ordering, atomic::workgroup);
+  switch (Ordering) {
+  default:
+    __builtin_unreachable();
+  case atomic::aquire:
+    return __builtin_amdgcn_fence(atomic::aquire, "workgroup");
+  case atomic::release:
+    return __builtin_amdgcn_fence(atomic::release, "workgroup");
+  case atomic::acq_rel:
+    return __builtin_amdgcn_fence(atomic::acq_rel, "workgroup");
+  case atomic::seq_cst:
+    return __builtin_amdgcn_fence(atomic::seq_cst, "workgroup");
+  }
 }
-
 void fenceKernel(atomic::OrderingTy Ordering) {
-  return __scoped_atomic_thread_fence(Ordering, atomic::device_);
+  switch (Ordering) {
+  default:
+    __builtin_unreachable();
+  case atomic::aquire:
+    return __builtin_amdgcn_fence(atomic::aquire, "agent");
+  case atomic::release:
+    return __builtin_amdgcn_fence(atomic::release, "agent");
+  case atomic::acq_rel:
+    return __builtin_amdgcn_fence(atomic::acq_rel, "agent");
+  case atomic::seq_cst:
+    return __builtin_amdgcn_fence(atomic::seq_cst, "agent");
+  }
 }
-
 void fenceSystem(atomic::OrderingTy Ordering) {
-  return __scoped_atomic_thread_fence(Ordering, atomic::system);
+  switch (Ordering) {
+  default:
+    __builtin_unreachable();
+  case atomic::aquire:
+    return __builtin_amdgcn_fence(atomic::aquire, "");
+  case atomic::release:
+    return __builtin_amdgcn_fence(atomic::release, "");
+  case atomic::acq_rel:
+    return __builtin_amdgcn_fence(atomic::acq_rel, "");
+  case atomic::seq_cst:
+    return __builtin_amdgcn_fence(atomic::seq_cst, "");
+  }
 }
 
 void syncWarp(__kmpc_impl_lanemask_t) {