[LIBCLC] Use amdgcn_fence builtin in mem_fence implementation (intel#10133)

Author: Jakub Chlanda

libclc/amdgcn-amdhsa/libspirv/synchronization/barrier.cl

@@ -8,44 +8,47 @@
 
 #include <clc/clc.h>
 #include <spirv/spirv.h>
+#include <spirv/spirv_types.h>
 
-void __clc_amdgcn_s_waitcnt(unsigned flags);
+#define BUILTIN_FENCE(semantics, scope_memory)                                 \
+  if (semantics & Acquire)                                                     \
+    return __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, scope_memory);             \
+  else if (semantics & Release)                                                \
+    return __builtin_amdgcn_fence(__ATOMIC_RELEASE, scope_memory);             \
+  else if (semantics & AcquireRelease)                                         \
+    return __builtin_amdgcn_fence(__ATOMIC_ACQ_REL, scope_memory);             \
+  else if (semantics & SequentiallyConsistent)                                 \
+    return __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, scope_memory);             \
+  else                                                                         \
+    return __builtin_amdgcn_fence(__ATOMIC_ACQ_REL, scope_memory);
 
-// s_waitcnt takes 16bit argument with a combined number of maximum allowed
-// pending operations:
-// [12:8] LGKM -- LDS, GDS, Konstant (SMRD), Messages
-// [7] -- undefined
-// [6:4] -- exports, GDS, and mem write
-// [3:0] -- vector memory operations
-
-// Newer clang supports __builtin_amdgcn_s_waitcnt
-#if __clang_major__ >= 5
-#define __waitcnt(x) __builtin_amdgcn_s_waitcnt(x)
-#else
-#define __waitcnt(x) __clc_amdgcn_s_waitcnt(x)
-_CLC_DEF void __clc_amdgcn_s_waitcnt(unsigned) __asm("llvm.amdgcn.s.waitcnt");
-#endif
-
-_CLC_DEF _CLC_OVERLOAD void __mem_fence(cl_mem_fence_flags flags) {
-  if (flags & CLK_GLOBAL_MEM_FENCE) {
-    // scalar loads are counted with LGKM but we don't know whether
-    // the compiler turned any loads to scalar
-    __waitcnt(0);
-  } else if (flags & CLK_LOCAL_MEM_FENCE)
-    __waitcnt(0xff); // LGKM is [12:8]
+_CLC_DEF _CLC_OVERLOAD void __mem_fence(unsigned int scope_memory,
+                                        unsigned int semantics) {
+  switch ((enum Scope)scope_memory) {
+  case CrossDevice:
+    BUILTIN_FENCE(semantics, "")
+  case Device:
+    BUILTIN_FENCE(semantics, "agent")
+  case Workgroup:
+    BUILTIN_FENCE(semantics, "workgroup")
+  case Subgroup:
+    BUILTIN_FENCE(semantics, "wavefront")
+  case Invocation:
+    BUILTIN_FENCE(semantics, "singlethread")
+  }
 }
-#undef __waitcnt
+#undef BUILTIN_FENCE
 
-_CLC_OVERLOAD _CLC_DEF void __spirv_MemoryBarrier(unsigned int memory,
+_CLC_OVERLOAD _CLC_DEF void __spirv_MemoryBarrier(unsigned int scope_memory,
                                                   unsigned int semantics) {
-  __mem_fence(memory);
+  __mem_fence(scope_memory, semantics);
 }
 
 _CLC_OVERLOAD _CLC_DEF _CLC_CONVERGENT void
-__spirv_ControlBarrier(unsigned int scope, unsigned int memory,
+__spirv_ControlBarrier(unsigned int scope_execution, unsigned scope_memory,
                        unsigned int semantics) {
   if (semantics) {
-    __mem_fence(memory);
+    __mem_fence(scope_memory, semantics);
   }
   __builtin_amdgcn_s_barrier();
 }