[SYCL][LIBCLC] Add Control/Memory Barrier builtins for AMDGCN

Author: jchlanda

libclc/amdgcn-amdhsa/libspirv/SOURCES

@@ -1,3 +1,4 @@
 workitem/get_global_size.cl
 workitem/get_local_size.cl
 workitem/get_num_groups.cl
+synchronization/barrier.cl

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

@@ -0,0 +1,50 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include <spirv/spirv.h>
+
+void __clc_amdgcn_s_waitcnt(unsigned flags);
+
+// 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]
+}
+#undef __waitcnt
+
+_CLC_OVERLOAD _CLC_DEF void __spirv_MemoryBarrier(unsigned int memory,
+                                                  unsigned int semantics) {
+  __mem_fence(memory);
+}
+
+_CLC_OVERLOAD _CLC_DEF _CLC_CONVERGENT void
+__spirv_ControlBarrier(unsigned int scope, unsigned int memory,
+                       unsigned int semantics) {
+  if (semantics) {
+    __mem_fence(memory);
+  }
+  __builtin_amdgcn_s_barrier();
+}

libclc/amdgcn/lib/mem_fence/fence.cl

@@ -1,37 +1,13 @@
 #include <clc/clc.h>
 
-void __clc_amdgcn_s_waitcnt(unsigned flags);
-
-// 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]
-}
-#undef __waitcnt
+_CLC_OVERLOAD _CLC_DEF void __spirv_MemoryBarrier(unsigned int memory,
+                                                  unsigned int semantics);
 
 // We don't have separate mechanism for read and write fences
 _CLC_DEF _CLC_OVERLOAD void read_mem_fence(cl_mem_fence_flags flags) {
-  mem_fence(flags);
+  __spirv_MemoryBarrier(flags, 1);
 }
 
 _CLC_DEF _CLC_OVERLOAD void write_mem_fence(cl_mem_fence_flags flags) {
-  mem_fence(flags);
+  __spirv_MemoryBarrier(flags, 1);
 }

libclc/amdgcn/lib/synchronization/barrier.cl

@@ -1,6 +1,11 @@
 #include <clc/clc.h>
 
+_CLC_OVERLOAD _CLC_DEF _CLC_CONVERGENT void
+__spirv_ControlBarrier(unsigned int scope, unsigned int memory,
+                       unsigned int semantics);
+
 _CLC_DEF _CLC_OVERLOAD void barrier(cl_mem_fence_flags flags) {
-  mem_fence(flags);
-  __builtin_amdgcn_s_barrier();
+  // Call spir-v implementation of the barrier.
+  // Set semantics to not None, so it performs mem fence.
+  __spirv_ControlBarrier(0, flags, 1);
 }