[libomptarget][nfc] Move omp locks under target_impl

[JonChesterfield]

Summary

[libomptarget][nfc] Move omp locks under target_impl

These are likely to be target specific, even down to the lock_t which is
correspondingly moved out of interface.h. The alternative is to include
interface.h in target_impl which substantiatially increases the scope of
those symbols.

The current nvptx implementation deadlocks on amdgcn. The preferred
implementation for that arch is still under discussion - this change
leaves declarations in target_impl.

The functions could be inline for nvptx. I'd prefer to keep the internals
hidden in the target_impl translation unit, but will add the (possibly renamed)
macros to target_impl.h if preferred.

[jdoerfert]

LGTM.

[JonChesterfield]

Thanks. This is most of the target specific code in libcall, the remaining function is for omp_get_wtime. I'll put up a patch for that shortly.

[JonChesterfield]

• Move print statements back into generic functions

[JonChesterfield]

Apologies for changing the patch post review. The print calls are now back in the omp_ prefixed functions. That seems a better separation of concerns.

openmp/libomptarget/deviceRTLs/amdgcn/src/amdgcn_interface.h

 //===--- amdgcn_interface.h - OpenMP interface definitions ------- CUDA -*-===//
 //
 // 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 _AMDGCN_INTERFACE_H_
 #define _AMDGCN_INTERFACE_H_
 
 #include <stdint.h>
 
 #define EXTERN extern "C" __attribute__((device))
 typedef uint64_t __kmpc_impl_lanemask_t;
+typedef uint32_t omp_lock_t; /* arbitrary type of the right length */
 
 #endif

openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h

 // Calls to the AMDGCN layer (assuming 1D layout)
 EXTERN uint64_t __ockl_get_local_size(uint32_t);
 EXTERN uint64_t __ockl_get_num_groups(uint32_t);
 INLINE int GetThreadIdInBlock() { return __builtin_amdgcn_workitem_id_x(); }
 INLINE int GetBlockIdInKernel() { return __builtin_amdgcn_workgroup_id_x(); }
 INLINE int GetNumberOfBlocksInKernel() { return __ockl_get_num_groups(0); }
 INLINE int GetNumberOfThreadsInBlock() { return __ockl_get_local_size(0); }
 
+// Locks
+EXTERN void __kmpc_impl_init_lock(omp_lock_t *lock);
+EXTERN void __kmpc_impl_destroy_lock(omp_lock_t *lock);
+EXTERN void __kmpc_impl_set_lock(omp_lock_t *lock);
+EXTERN void __kmpc_impl_unset_lock(omp_lock_t *lock);
+EXTERN int __kmpc_impl_test_lock(omp_lock_t *lock);
+
 // DEVICE versions of part of libc
 extern "C" {
 DEVICE __attribute__((noreturn)) void
 __assertfail(const char *, const char *, unsigned, const char *, size_t);
 INLINE static void __assert_fail(const char *__message, const char *__file,
                                  unsigned int __line, const char *__function) {
   __assertfail(__message, __file, __line, __function, sizeof(char));
 }
 DEVICE int printf(const char *, ...);
 }
 
 #endif

openmp/libomptarget/deviceRTLs/interface.h

 #ifdef __CUDACC__
 #include "nvptx/src/nvptx_interface.h"
 #endif
 
 ////////////////////////////////////////////////////////////////////////////////
 // OpenMP interface
 ////////////////////////////////////////////////////////////////////////////////
 
-typedef uint32_t omp_lock_t;      /* arbitrary type of the right length */
 typedef uint64_t omp_nest_lock_t; /* arbitrary type of the right length */
 
 typedef enum omp_sched_t {
   omp_sched_static = 1,  /* chunkSize >0 */
   omp_sched_dynamic = 2, /* chunkSize >0 */
   omp_sched_guided = 3,  /* chunkSize >0 */
   omp_sched_auto = 4,    /* no chunkSize */
 } omp_sched_t;

openmp/libomptarget/deviceRTLs/nvptx/CMakeLists.txt

if(LIBOMPTARGET_DEP_CUDA_FOUND)
   # propagating host flags.
   set(CUDA_PROPAGATE_HOST_FLAGS OFF)
 
   set(cuda_src_files
       ${devicertl_common_directory}/src/cancel.cu
       ${devicertl_common_directory}/src/critical.cu
       src/data_sharing.cu
       src/libcall.cu
+      src/target_impl.cu
       ${devicertl_common_directory}/src/loop.cu
       ${devicertl_common_directory}/src/omptarget.cu
       ${devicertl_common_directory}/src/parallel.cu
       src/reduction.cu
       ${devicertl_common_directory}/src/support.cu
       ${devicertl_common_directory}/src/sync.cu
       ${devicertl_common_directory}/src/task.cu
   )

openmp/libomptarget/deviceRTLs/nvptx/src/libcall.cu

EXTERN int omp_get_max_task_priority(void) {
   PRINT0(LD_IO, "call omp_get_max_task_priority() returns 0\n");
   return 0;
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 // locks
 ////////////////////////////////////////////////////////////////////////////////
 
-#define __OMP_SPIN 1000
-#define UNSET 0
-#define SET 1
-
 EXTERN void omp_init_lock(omp_lock_t *lock) {
-  omp_unset_lock(lock);
+  __kmpc_impl_init_lock(lock);
   PRINT0(LD_IO, "call omp_init_lock()\n");
 }
 
 EXTERN void omp_destroy_lock(omp_lock_t *lock) {
-  omp_unset_lock(lock);
+  __kmpc_impl_destroy_lock(lock);
   PRINT0(LD_IO, "call omp_destroy_lock()\n");
 }
 
 EXTERN void omp_set_lock(omp_lock_t *lock) {
-  // int atomicCAS(int* address, int compare, int val);
-  // (old == compare ? val : old)
-
-  // TODO: not sure spinning is a good idea here..
-  while (atomicCAS(lock, UNSET, SET) != UNSET) {
-    clock_t start = clock();
-    clock_t now;
-    for (;;) {
-      now = clock();
-      clock_t cycles = now > start ? now - start : now + (0xffffffff - start);
-      if (cycles >= __OMP_SPIN * GetBlockIdInKernel()) {
-        break;
-      }
-    }
-  } // wait for 0 to be the read value
-
+  __kmpc_impl_set_lock(lock);
   PRINT0(LD_IO, "call omp_set_lock()\n");
 }
 
 EXTERN void omp_unset_lock(omp_lock_t *lock) {
-  (void)atomicExch(lock, UNSET);
-
+  __kmpc_impl_unset_lock(lock);
   PRINT0(LD_IO, "call omp_unset_lock()\n");
 }
 
 EXTERN int omp_test_lock(omp_lock_t *lock) {
-  // int atomicCAS(int* address, int compare, int val);
-  // (old == compare ? val : old)
-  int ret = atomicAdd(lock, 0);
-
+  int rc = __kmpc_impl_test_lock(lock);
   PRINT(LD_IO, "call omp_test_lock() return %d\n", ret);
-
-  return ret;
+  return rc;
 }
 
 // for xlf Fotran
 // Fotran, the return is LOGICAL type
 
 #define FLOGICAL long
 EXTERN FLOGICAL __xlf_omp_is_initial_device_i8() {
   int ret = omp_is_initial_device();

openmp/libomptarget/deviceRTLs/nvptx/src/nvptx_interface.h

 //===--- nvptx_interface.h - OpenMP interface definitions -------- CUDA -*-===//
 //
 // 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 _NVPTX_INTERFACE_H_
 #define _NVPTX_INTERFACE_H_
 
 #include <stdint.h>
 
 #define EXTERN extern "C" __device__
 typedef uint32_t __kmpc_impl_lanemask_t;
+typedef uint32_t omp_lock_t; /* arbitrary type of the right length */
 
 #endif

openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h

 INLINE void __kmpc_impl_threadfence_system(void) { __threadfence_system(); }
 
 // Calls to the NVPTX layer (assuming 1D layout)
 INLINE int GetThreadIdInBlock() { return threadIdx.x; }
 INLINE int GetBlockIdInKernel() { return blockIdx.x; }
 INLINE int GetNumberOfBlocksInKernel() { return gridDim.x; }
 INLINE int GetNumberOfThreadsInBlock() { return blockDim.x; }
 
+// Locks
+EXTERN void __kmpc_impl_init_lock(omp_lock_t *lock);
+EXTERN void __kmpc_impl_destroy_lock(omp_lock_t *lock);
+EXTERN void __kmpc_impl_set_lock(omp_lock_t *lock);
+EXTERN void __kmpc_impl_unset_lock(omp_lock_t *lock);
+EXTERN int __kmpc_impl_test_lock(omp_lock_t *lock);
+
 #endif

openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.cu

+//===---------- target_impl.cu - NVPTX OpenMP GPU options ------- CUDA -*-===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// Definitions of target specific functions
+//
+//===----------------------------------------------------------------------===//
+
+#include "target_impl.h"
+#include "common/debug.h"
+
+#define __OMP_SPIN 1000
+#define UNSET 0
+#define SET 1
+
+EXTERN void __kmpc_impl_init_lock(omp_lock_t *lock) {
+  omp_unset_lock(lock);
+}
+
+EXTERN void __kmpc_impl_destroy_lock(omp_lock_t *lock) {
+  omp_unset_lock(lock);
+}
+
+EXTERN void __kmpc_impl_set_lock(omp_lock_t *lock) {
+  // int atomicCAS(int* address, int compare, int val);
+  // (old == compare ? val : old)
+
+  // TODO: not sure spinning is a good idea here..
+  while (atomicCAS(lock, UNSET, SET) != UNSET) {
+    clock_t start = clock();
+    clock_t now;
+    for (;;) {
+      now = clock();
+      clock_t cycles = now > start ? now - start : now + (0xffffffff - start);
+      if (cycles >= __OMP_SPIN * GetBlockIdInKernel()) {
+        break;
+      }
+    }
+  } // wait for 0 to be the read value
+}
+
+EXTERN void __kmpc_impl_unset_lock(omp_lock_t *lock) {
+  (void)atomicExch(lock, UNSET);
+}
+
+EXTERN int __kmpc_impl_test_lock(omp_lock_t *lock) {
+  // int atomicCAS(int* address, int compare, int val);
+  // (old == compare ? val : old)
+  return atomicAdd(lock, 0);
+}