[libclc] Move mem_fence and barrier to clc library

libclc/clc/include/clc/mem_fence/clc_mem_fence.h

@@ -0,0 +1,17 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CLC_MEM_FENCE_CLC_MEM_FENCE_H__
+#define __CLC_MEM_FENCE_CLC_MEM_FENCE_H__
+
+#include <clc/internal/clc.h>
+
+_CLC_OVERLOAD _CLC_DECL void __clc_mem_fence(int memory_scope,
+                                             int memory_order);
+
+#endif // __CLC_MEM_FENCE_CLC_MEM_FENCE_H__

libclc/clc/include/clc/synchronization/clc_work_group_barrier.h

@@ -0,0 +1,17 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CLC_SYNCHRONIZATION_CLC_WORK_GROUP_BARRIER_H__
+#define __CLC_SYNCHRONIZATION_CLC_WORK_GROUP_BARRIER_H__
+
+#include <clc/internal/clc.h>
+
+_CLC_OVERLOAD _CLC_DECL void __clc_work_group_barrier(int memory_scope,
+                                                      int memory_order);
+
+#endif // __CLC_SYNCHRONIZATION_CLC_WORK_GROUP_BARRIER_H__

libclc/clc/lib/amdgcn/SOURCES

@@ -1,6 +1,8 @@
 math/clc_fmax.cl
 math/clc_fmin.cl
 math/clc_ldexp_override.cl
+mem_fence/clc_mem_fence.cl
+synchronization/clc_work_group_barrier.cl
 workitem/clc_get_global_offset.cl
 workitem/clc_get_global_size.cl
 workitem/clc_get_group_id.cl

libclc/clc/lib/amdgcn/mem_fence/clc_mem_fence.cl

@@ -0,0 +1,37 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 <clc/mem_fence/clc_mem_fence.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_OVERLOAD _CLC_DEF void __clc_mem_fence(int memory_scope,
+                                            int memory_order) {
+  if (memory_scope & __MEMORY_SCOPE_DEVICE) {
+    // scalar loads are counted with LGKM but we don't know whether
+    // the compiler turned any loads to scalar
+    __waitcnt(0);
+  } else if (memory_scope & __MEMORY_SCOPE_WRKGRP)
+    __waitcnt(0xff); // LGKM is [12:8]
+}
+#undef __waitcnt

libclc/clc/lib/amdgcn/synchronization/clc_work_group_barrier.cl

@@ -0,0 +1,16 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 <clc/mem_fence/clc_mem_fence.h>
+#include <clc/synchronization/clc_work_group_barrier.h>
+
+_CLC_OVERLOAD _CLC_DEF void __clc_work_group_barrier(int memory_scope,
+                                                     int memory_order) {
+  __clc_mem_fence(memory_scope, memory_order);
+  __builtin_amdgcn_s_barrier();
+}

libclc/clc/lib/ptx-nvidiacl/SOURCES

@@ -1,3 +1,5 @@
+mem_fence/clc_mem_fence.cl
+synchronization/clc_work_group_barrier.cl
 workitem/clc_get_global_id.cl
 workitem/clc_get_group_id.cl
 workitem/clc_get_local_id.cl

libclc/clc/lib/ptx-nvidiacl/mem_fence/clc_mem_fence.cl

@@ -0,0 +1,15 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 <clc/mem_fence/clc_mem_fence.h>
+
+_CLC_OVERLOAD _CLC_DEF void __clc_mem_fence(int memory_scope,
+                                            int memory_order) {
+  if (memory_scope & (__MEMORY_SCOPE_DEVICE | __MEMORY_SCOPE_WRKGRP))
+    __nvvm_membar_cta();
+}

libclc/clc/lib/ptx-nvidiacl/synchronization/clc_work_group_barrier.cl

@@ -0,0 +1,14 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 <clc/synchronization/clc_work_group_barrier.h>
+
+_CLC_OVERLOAD _CLC_DEF void __clc_work_group_barrier(int memory_scope,
+                                                     int memory_order) {
+  __syncthreads();
+}

libclc/opencl/include/clc/opencl/synchronization/cl_mem_fence_flags.h

@@ -13,5 +13,6 @@ typedef uint cl_mem_fence_flags;
 
 #define CLK_LOCAL_MEM_FENCE 1
 #define CLK_GLOBAL_MEM_FENCE 2
+#define CLK_IMAGE_MEM_FENCE 4
 
 #endif // __CLC_OPENCL_SYNCHRONIZATION_CL_MEM_FENCE_FLAGS_H__

libclc/opencl/include/clc/opencl/synchronization/utils.h

@@ -0,0 +1,24 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CLC_OPENCL_SYNCHRONIZATION_UTILS_H__
+#define __CLC_OPENCL_SYNCHRONIZATION_UTILS_H__
+
+#include <clc/internal/clc.h>
+#include <clc/opencl/synchronization/cl_mem_fence_flags.h>
+
+_CLC_INLINE int getCLCMemoryScope(cl_mem_fence_flags flag) {
+  int memory_scope = 0;
+  if (flag & CLK_GLOBAL_MEM_FENCE)
+    memory_scope |= __MEMORY_SCOPE_DEVICE;
+  if (flag & CLK_LOCAL_MEM_FENCE)
+    memory_scope |= __MEMORY_SCOPE_WRKGRP;
+  return memory_scope;
+}
+
+#endif // __CLC_OPENCL_SYNCHRONIZATION_UTILS_H__

libclc/opencl/lib/amdgcn/mem_fence/fence.cl

@@ -6,34 +6,15 @@
 //
 //===----------------------------------------------------------------------===//
 
+#include <clc/mem_fence/clc_mem_fence.h>
 #include <clc/opencl/explicit_fence/explicit_memory_fence.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
+#include <clc/opencl/synchronization/utils.h>
 
 _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]
+  int memory_scope = getCLCMemoryScope(flags);
+  int memory_order = __ATOMIC_SEQ_CST;
+  __clc_mem_fence(memory_scope, memory_order);
 }
-#undef __waitcnt
 
 // We don't have separate mechanism for read and write fences
 _CLC_DEF _CLC_OVERLOAD void read_mem_fence(cl_mem_fence_flags flags) {

libclc/opencl/lib/amdgcn/synchronization/barrier.cl

@@ -6,10 +6,12 @@
 //
 //===----------------------------------------------------------------------===//
 
-#include <clc/opencl/explicit_fence/explicit_memory_fence.h>
 #include <clc/opencl/synchronization/barrier.h>
+#include <clc/opencl/synchronization/utils.h>
+#include <clc/synchronization/clc_work_group_barrier.h>
 
 _CLC_DEF _CLC_OVERLOAD void barrier(cl_mem_fence_flags flags) {
-  mem_fence(flags);
-  __builtin_amdgcn_s_barrier();
+  int memory_scope = getCLCMemoryScope(flags);
+  int memory_order = __ATOMIC_SEQ_CST;
+  __clc_work_group_barrier(memory_scope, memory_order);
 }

libclc/opencl/lib/ptx-nvidiacl/mem_fence/fence.cl

@@ -6,11 +6,14 @@
 //
 //===----------------------------------------------------------------------===//
 
+#include <clc/mem_fence/clc_mem_fence.h>
 #include <clc/opencl/explicit_fence/explicit_memory_fence.h>
+#include <clc/opencl/synchronization/utils.h>
 
 _CLC_DEF _CLC_OVERLOAD void mem_fence(cl_mem_fence_flags flags) {
-  if (flags & (CLK_GLOBAL_MEM_FENCE | CLK_LOCAL_MEM_FENCE))
-    __nvvm_membar_cta();
+  int memory_scope = getCLCMemoryScope(flags);
+  int memory_order = __ATOMIC_SEQ_CST;
+  __clc_mem_fence(memory_scope, memory_order);
 }
 
 // We do not have separate mechanism for read and write fences.

libclc/opencl/lib/ptx-nvidiacl/synchronization/barrier.cl

@@ -7,7 +7,11 @@
 //===----------------------------------------------------------------------===//
 
 #include <clc/opencl/synchronization/barrier.h>
+#include <clc/opencl/synchronization/utils.h>
+#include <clc/synchronization/clc_work_group_barrier.h>
 
 _CLC_DEF _CLC_OVERLOAD void barrier(cl_mem_fence_flags flags) {
-  __syncthreads();
+  int memory_scope = getCLCMemoryScope(flags);
+  int memory_order = __ATOMIC_SEQ_CST;
+  __clc_work_group_barrier(memory_scope, memory_order);
 }