[SYCL][LIBCLC] Add memory/control barrier, sqrt, sin and cos builtins for AMDGCN (#4091)

Author: jchlanda

libclc/CMakeLists.txt

@@ -343,6 +343,10 @@ foreach( t ${LIBCLC_TARGETS_TO_BUILD} )
 			set( opt_flags -O3 )
 		endif()
 
+                # Enable SPIR-V builtin function declarations, so they don't
+                # have to be explicity declared in the soruce.
+                list( APPEND flags -Xclang -fdeclare-spirv-builtins)
+
 		add_libclc_builtin_set(libspirv-${arch_suffix}
 			TRIPLE ${t}
 			TARGET_ENV libspirv

libclc/amdgcn-amdhsa/libspirv/SOURCES

@@ -1,3 +1,7 @@
 workitem/get_global_size.cl
 workitem/get_local_size.cl
 workitem/get_num_groups.cl
+synchronization/barrier.cl
+math/cos.cl
+math/sin.cl
+math/sqrt.cl

libclc/amdgcn-amdhsa/libspirv/math/cos.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 <spirv/spirv.h>
+
+_CLC_OVERLOAD _CLC_DECL _CLC_CONSTFN __clc_fp32_t
+__spirv_ocl_cos(__clc_fp32_t In) {
+  return __builtin_amdgcn_cosf(In);
+}

libclc/amdgcn-amdhsa/libspirv/math/sin.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 <spirv/spirv.h>
+
+_CLC_OVERLOAD _CLC_DECL _CLC_CONSTFN __clc_fp32_t
+__spirv_ocl_sin(__clc_fp32_t In) {
+  return __builtin_amdgcn_sinf(In);
+}

libclc/amdgcn-amdhsa/libspirv/math/sqrt.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 <spirv/spirv.h>
+
+_CLC_OVERLOAD _CLC_DECL _CLC_CONSTFN __clc_fp32_t
+__spirv_ocl_sqrt(__clc_fp32_t In) {
+  return __builtin_amdgcn_sqrtf(In);
+}

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,10 @@
 #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
-
 // 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,7 @@
 #include <clc/clc.h>
 
 _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);
 }