[OpenMP][WIP][POC] Compile the device runtime as C++

[jdoerfert]

Summary

The main contribution of this patch is the new CMakeLists.txt which
shows how we can compile the device runtime as C++ for various targets
and subtargets. Running ninja omptarget-device-rtl results in

libomptarget-device-rtl-nvptx64-sm35-ptx50.bc
libomptarget-device-rtl-nvptx64-sm35-ptx60.bc
libomptarget-device-rtl-nvptx64-sm35-ptx61.bc
libomptarget-device-rtl-nvptx64-sm35-ptx62.bc
libomptarget-device-rtl-nvptx64-sm35-ptx63.bc
libomptarget-device-rtl-nvptx64-sm35-ptx64.bc
libomptarget-device-rtl-nvptx64-sm35-ptx65.bc
libomptarget-device-rtl-nvptx64-sm35-ptx70.bc
libomptarget-device-rtl-nvptx64-sm35-ptx71.bc
libomptarget-device-rtl-nvptx-sm35-ptx50.bc
libomptarget-device-rtl-nvptx-sm35-ptx60.bc
libomptarget-device-rtl-nvptx-sm35-ptx61.bc
libomptarget-device-rtl-nvptx-sm35-ptx62.bc
libomptarget-device-rtl-nvptx-sm35-ptx63.bc
libomptarget-device-rtl-nvptx-sm35-ptx64.bc
libomptarget-device-rtl-nvptx-sm35-ptx65.bc
libomptarget-device-rtl-nvptx-sm35-ptx70.bc
libomptarget-device-rtl-nvptx-sm35-ptx71.bc
libomptarget-device-rtl-x86_64-unknown-linux-gnu.bc

As part of this we should remove most of the source annotations, e.g.
INLINE, and EXTERN, as long as they are not strictly needed. SHARED
will be replaced by the macro used in D93135 or D77472.

For lanemask and similar target dependent functions we should have a
begin/end declare variant file:

// I'm gussing here:
using kmpc_lanemask = uint64_t;

void some_target_specific_function_decl(...);

void __kmpc_some_target_specific_function_impl(...) {
  some_target_specific_function_decl(...);
}

[jdoerfert]

Remove accidentially included header

[jdoerfert]

openmp/libomptarget/deviceRTLs/common/debug.h

129

All of this seems totally broken. As far as I can tell, OMPTARGET_NVPTX_TEST and some other macros are always set and later we check if they are set. It's odd.

openmp/libomptarget/deviceRTLs/common/omptarget.h

41

If we really want, we can mark the entry points with the flatten attribute. Other than that, there is little point in adding "inline" and such everywhere through these macros. If the compiler doesn't figure it out we should change the target pipeline defaults instead and thereby help user code as well.

openmp/libomptarget/deviceRTLs/common/src/cancel.cu

15

We should go through and minimize include dependences before we do anything else. The structure right now is also very confusing. There is no need for "internal headers" and the likes anymore.

openmp/libomptarget/deviceRTLs/CMakeLists.txt

 ##===----------------------------------------------------------------------===##
 #
 # 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
 #
 # ##===----------------------------------------------------------------------===##
 #
 # Build a device RTL for each available machine.
 #
 ##===----------------------------------------------------------------------===##
 
+add_subdirectory(common)
 add_subdirectory(nvptx)

openmp/libomptarget/deviceRTLs/common/CMakeLists.txt

+##===----------------------------------------------------------------------===##
+#
+# 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
+#
+##===----------------------------------------------------------------------===##
+#
+# Build the Common Device RTL parts.
+#
+##===----------------------------------------------------------------------===##
+
+set(LIBOMPTARGET_DEVICE_RTL_COMPILER "" CACHE STRING
+  "Path to alternate compiler to be used for the device RTL.")
+set(LIBOMPTARGET_DEVICE_RTL_LINKER "" CACHE STRING
+  "Path to alternate compiler to be used for the device RTL.")
+
+if(LIBOMPTARGET_DEVICE_RTL_COMPILER)
+  find_program(ALTERNATE_DEVICE_RTL_COMPILER NAMES ${LIBOMPTARGET_DEVICE_RTL_COMPILER})
+  if(NOT ALTERNATE_DEVICE_RTL_COMPILER)
+    libomptarget_say("Not building offloading device RTL:
+    LIBOMPTARGET_DEVICE_RTL_COMPILER not valid, was '${LIBOMPTARGET_DEVICE_RTL_LINKER}'")
+  endif()
+  set(LIBOMPTARGET_DEVICE_RTL_COMPILER ${ALTERNATE_DEVICE_RTL_COMPILER} CACHE FILEPATH "" FORCE)
+else()
+  set(LIBOMPTARGET_DEVICE_RTL_COMPILER ${CMAKE_CXX_COMPILER} CACHE FILEPATH "" FORCE)
+endif()
+
+if(NOT LIBOMPTARGET_DEVICE_RTL_LINKER)
+  get_filename_component(compiler_base_directory ${LIBOMPTARGET_DEVICE_RTL_COMPILER} DIRECTORY)
+  set(LIBOMPTARGET_DEVICE_RTL_LINKER "${compiler_base_directory}/llvm-link" CACHE FILEPATH "" FORCE)
+endif()
+
+if(LIBOMPTARGET_DEVICE_RTL_LINKER)
+  find_program(ALTERNATE_DEVICE_RTL_LINKER NAMES ${LIBOMPTARGET_DEVICE_RTL_LINKER})
+  if(NOT ALTERNATE_DEVICE_RTL_LINKER)
+    libomptarget_say("Not building offloading device RTL:
+    LIBOMPTARGET_DEVICE_RTL_LINKER not valid, was '${LIBOMPTARGET_DEVICE_RTL_LINKER}'")
+  endif()
+  set(LIBOMPTARGET_DEVICE_RTL_LINKER ${ALTERNATE_DEVICE_RTL_LINKER} CACHE FILEPATH "" FORCE)
+endif()
+
+if(LIBOMPTARGET_DEVICE_RTL_COMPILER AND LIBOMPTARGET_DEVICE_RTL_LINKER)
+
+  get_filename_component(devicertl_base_directory
+    ${CMAKE_CURRENT_SOURCE_DIR}
+    DIRECTORY)
+
+  set(src_files
+      src/cancel.cu
+      #src/loop.cu
+  )
+
+
+  libomptarget_say("Building LLVM bitcode offloading device RTL.")
+
+  # Set flags for LLVM Bitcode compilation.
+  set(bc_flags  -S
+                -x c++
+                -fopenmp
+                -Xclang -fopenmp-is-device
+                -Xclang -emit-llvm-bc
+                -I${devicertl_base_directory}
+                -I${CMAKE_CURRENT_SOURCE_DIR}
+                -I${CMAKE_CURRENT_SOURCE_DIR}/include
+              )
+
+  # Activate RTL message dumps if requested by the user.
+  set(LIBOMPTARGET_DEBUG FALSE CACHE BOOL
+      "Activate NVPTX device RTL debug messages.")
+
+  if(${LIBOMPTARGET_DEBUG})
+    set(bc_flags ${bc_flags} -DOMPTARGET_DEBUG=-1)
+  else()
+    set(bc_flags ${bc_flags} -DOMPTARGET_DEBUG=0)
+  endif()
+
+  # Create target to build all Bitcode libraries.
+  add_custom_target(omptarget-device-rtl)
+
+  set(names "")
+  set(triples "")
+  set(target_cpus "")
+  set(target_features "")
+  set(aux_triples "")
+
+  # NVIDIA configurations
+
+  # Generate a Bitcode library for all the compute capabilities the user requested.
+  set(nvptx_ptx_list 50 60 61 62 63 64 65 70 71)
+  set(nvptx_triple_list "nvptx" "nvptx64")
+  #execute_process(COMMAND ${LIBOMPTARGET_DEVICE_RTL_COMPILER} -v | grep 'target' OUTPUT_VARIABLE aux_triple)
+  set(nvptx_aux_triples "x86_64-unknown-linux-gnu")
+
+  # Build library support for the highest compute capability the system supports
+  # and always build support for sm_35 by default
+  if (${LIBOMPTARGET_DEP_CUDA_ARCH} EQUAL 35)
+      set(default_capabilities 35)
+  else()
+      set(default_capabilities "35,${LIBOMPTARGET_DEP_CUDA_ARCH}")
+  endif()
+
+  if (DEFINED LIBOMPTARGET_NVPTX_COMPUTE_CAPABILITY)
+    set(default_capabilities ${LIBOMPTARGET_NVPTX_COMPUTE_CAPABILITY})
+    libomptarget_warning_say("LIBOMPTARGET_NVPTX_COMPUTE_CAPABILITY is deprecated, please use LIBOMPTARGET_NVPTX_COMPUTE_CAPABILITIES")
+  endif()
+  set(LIBOMPTARGET_NVPTX_COMPUTE_CAPABILITIES ${default_capabilities} CACHE STRING
+    "List of CUDA Compute Capabilities to be used to compile the NVPTX device RTL.")
+  string(REPLACE "," ";" nvptx_sm_list ${LIBOMPTARGET_NVPTX_COMPUTE_CAPABILITIES})
+
+
+  foreach(ptx ${nvptx_ptx_list})
+    foreach(triple ${nvptx_triple_list})
+      foreach(sm ${nvptx_sm_list})
+      foreach(aux_triple ${nvptx_aux_triples})
+          list(APPEND names "${triple}-sm${sm}-ptx${ptx}")
+          list(APPEND triples ${triple})
+          list(APPEND target_cpus sm_${sm})
+          list(APPEND target_features +ptx${ptx})
+          list(APPEND aux_triples ${aux_triple})
+      endforeach()
+      endforeach()
+    endforeach()
+  endforeach()
+
+  # - NVIDIA
+
+  # Host
+
+  set(host_triple "x86_64-unknown-linux-gnu")
+  list(APPEND names "${host_triple}")
+  list(APPEND triples ${host_triple})
+  list(APPEND target_cpus "")
+  list(APPEND target_features "")
+  list(APPEND aux_triples "")
+
+  # - Host
+
+
+  foreach(cname triple target_cpu target_feature aux_triple IN ZIP_LISTS names triples target_cpus target_features aux_triples)
+    # Compile CUDA files to bitcode.
+    libomptarget_say("${cname} --- ${triple} --- ${target_cpu}")
+    set(bc_files "")
+    foreach(src ${src_files})
+      get_filename_component(infile ${src} ABSOLUTE)
+      get_filename_component(outfile ${src} NAME)
+
+      set(config_flags "")
+      if(${triple})
+        list(APPEND config_flags -target=${triple})
+      endif()
+      if(${target_cpu})
+        list(APPEND config_flags -mcpu=${target_cpu})
+      endif()
+      if(${target_feature})
+        list(APPEND config_flags -Xclang)
+        list(APPEND config_flags -target-feature)
+        list(APPEND config_flags -Xclang)
+        list(APPEND config_flags ${target_feature})
+      endif()
+      if(${aux_triple})
+        list(APPEND config_flags -Xclang)
+        list(APPEND config_flags -aux-triple)
+        list(APPEND config_flags -Xclang)
+        list(APPEND config_flags ${aux_triple})
+      endif()
+
+      add_custom_command(OUTPUT ${outfile}-${cname}.bc
+      COMMAND ${LIBOMPTARGET_DEVICE_RTL_COMPILER} ${bc_flags}
+                ${config_flags}
+                ${} ${infile} -o ${outfile}-${cname}.bc
+        DEPENDS ${infile}
+        IMPLICIT_DEPENDS CXX ${infile}
+        COMMENT "Building LLVM bitcode ${outfile}-${cname}.bc"
+        VERBATIM
+      )
+      set_property(DIRECTORY APPEND PROPERTY ADDITIONAL_MAKE_CLEAN_FILES
+      ${outfile}-${cname}.bc)
+
+      list(APPEND bc_files ${outfile}-${cname}.bc)
+    endforeach()
+
+    # Link to a bitcode library.
+    add_custom_command(OUTPUT
+    ${CMAKE_CURRENT_BINARY_DIR}/libomptarget-device-rtl-${cname}.bc
+      COMMAND ${LIBOMPTARGET_DEVICE_RTL_LINKER}
+          -o ${CMAKE_CURRENT_BINARY_DIR}/libomptarget-device-rtl-${cname}.bc ${bc_files}
+        DEPENDS ${bc_files}
+        COMMENT "Linking LLVM bitcode libomptarget-device-rtl-${cname}.bc"
+    )
+    set_property(DIRECTORY APPEND PROPERTY ADDITIONAL_MAKE_CLEAN_FILES
+    libomptarget-device-rtl-${cname}.bc)
+
+    add_custom_target(omptarget-device-rtl-${cname}-bc ALL DEPENDS
+    ${CMAKE_CURRENT_BINARY_DIR}/libomptarget-device-rtl-${cname}.bc)
+    add_dependencies(omptarget-device-rtl omptarget-device-rtl-${cname}-bc)
+
+    # Copy library to destination.
+    add_custom_command(TARGET omptarget-device-rtl-${cname}-bc POST_BUILD
+                        COMMAND ${CMAKE_COMMAND} -E copy
+                        ${CMAKE_CURRENT_BINARY_DIR}/libomptarget-device-rtl-${cname}.bc
+                        ${OPENMP_INSTALL_LIBDIR})
+
+    # Install bitcode library under the lib destination folder.
+    install(FILES
+    ${CMAKE_CURRENT_BINARY_DIR}/libomptarget-device-rtl-${cname}.bc DESTINATION "${OPENMP_INSTALL_LIBDIR}")
+  endforeach()
+endif()

openmp/libomptarget/deviceRTLs/common/debug.h

 //   WARNING(flag, str, args)      : if warning flag is on, print the warning
 //   WON(flag)                     : return true if warning flag is on
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef _OMPTARGET_NVPTX_DEBUG_H_
 #define _OMPTARGET_NVPTX_DEBUG_H_
 
 #include "common/device_environment.h"
 
 ////////////////////////////////////////////////////////////////////////////////
 // set desired level of debugging
 ////////////////////////////////////////////////////////////////////////////////
 
 #define LD_SET_NONE 0ULL /* none */
 #define LD_SET_ALL -1ULL /* all */
 

 #else
 #define OMPTARGET_NVPTX_WARNING (LW_SET_FUSSY)
 #endif
 
 ////////////////////////////////////////////////////////////////////////////////
 // implementation for debug
 ////////////////////////////////////////////////////////////////////////////////
 
-#if OMPTARGET_NVPTX_DEBUG || OMPTARGET_NVPTX_TEST || OMPTARGET_NVPTX_WARNING
+#if OMPTARGET_NVPTX_DEBUG
 #include "common/support.h"
 
 template <typename... Arguments>
 NOINLINE static void log(const char *fmt, Arguments... parameters) {
   printf(fmt, (int)GetBlockIdInKernel(), (int)GetThreadIdInBlock(),
          (int)GetWarpId(), (int)GetLaneId(), parameters...);
 }
 
-#endif
-#if OMPTARGET_NVPTX_TEST
-
 template <typename... Arguments>
 NOINLINE static void check(bool cond, const char *fmt,
                            Arguments... parameters) {
   if (!cond)
     printf(fmt, (int)GetBlockIdInKernel(), (int)GetThreadIdInBlock(),
            (int)GetWarpId(), (int)GetLaneId(), parameters...);
   assert(cond);
 }

 ////////////////////////////////////////////////////////////////////////////////
 // early defs for test
 ////////////////////////////////////////////////////////////////////////////////
 
 #define LT_SAFETY (LT_SET_SAFETY | LT_SET_INPUT | LT_SET_FUSSY)
 #define LT_INPUT (LT_SET_INPUT | LT_SET_FUSSY)
 #define LT_FUSSY (LT_SET_FUSSY)
 
+#if OMPTARGET_NVPTX_DEBUG
 #if OMPTARGET_NVPTX_TEST == LT_SET_SAFETY
 
 #define TON(_flag) ((OMPTARGET_NVPTX_TEST) & (_flag))
 #define ASSERT0(_flag, _cond, _str)                                            \
   {                                                                            \
     if (TON(_flag)) {                                                          \
       check(_cond);                                                            \
     }                                                                          \

 #define ASSERT(_flag, _cond, _str, _args...)                                   \
   {                                                                            \
     if (TON(_flag)) {                                                          \
       check((_cond), "<b %3d, t %4d, w %2d, l %d2> ASSERT: " _str "\n",        \
             _args);                                                            \
     }                                                                          \
   }
 
+#endif
 #else
 
 #define TON(_flag) (0)
 #define ASSERT0(_flag, _cond, _str)
 #define ASSERT(_flag, _cond, _str, _args...)
 
 #endif
 
 ////////////////////////////////////////////////////////////////////////////////
 // early defs for warning
 
 #define LW_ALL (LW_SET_ALL)
 #define LW_ENV (LW_SET_FUSSY | LW_SET_INPUT | LW_SET_ENV)
 #define LW_INPUT (LW_SET_FUSSY | LW_SET_INPUT)
 #define LW_FUSSY (LW_SET_FUSSY)
 
-#if OMPTARGET_NVPTX_WARNING
+#if OMPTARGET_NVPTX_DEBUG
 
 #define WON(_flag) ((OMPTARGET_NVPTX_WARNING) & (_flag))
 #define WARNING0(_flag, _str)                                                  \
   {                                                                            \
     if (WON(_flag)) {                                                          \
       log("<b %2d, t %4d, w %2d, l %2d> WARNING: " _str);                      \
     }                                                                          \
   }

openmp/libomptarget/deviceRTLs/common/device_environment.h

 //===---- device_environment.h - OpenMP GPU device environment --- 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
 //
 //===----------------------------------------------------------------------===//
 //
 // Global device environment
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef _OMPTARGET_DEVICE_ENVIRONMENT_H_
 #define _OMPTARGET_DEVICE_ENVIRONMENT_H_
 
-#include "target_impl.h"
+#include <stdint.h>
 
 struct omptarget_device_environmentTy {
   int32_t debug_level;
 };
 
-extern DEVICE omptarget_device_environmentTy omptarget_device_environment;
+extern omptarget_device_environmentTy omptarget_device_environment;
 
 #endif

openmp/libomptarget/deviceRTLs/common/include/target.h

+//===-- target.cpp -------- OpenMP device runtime target implementation ---===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// Target region interfaces are simple interfaces designed to allow middle-end
+// (=LLVM) passes to analyze and transform the code. To achieve good performance
+// it may be required to run the associated passes. However, implementations of
+// this interface shall always provide a correct implementation as close to the
+// user expected code as possible.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_OPENMP_LIBOMPTARGET_DEVICERTLS_COMMON_TARGET_H
+#define LLVM_OPENMP_LIBOMPTARGET_DEVICERTLS_COMMON_TARGET_H
+
+#include <stdint.h>
+
+#ifndef EXTERN
+#define EXTERN
+#endif
+#ifndef CALLBACK
+#define CALLBACK(Callee, Payload0, Payload1)
+#endif
+
+/// Forward declaration of the source location identifier "ident".
+typedef struct ident ident_t;
+
+/// The target region _kernel_ interface for GPUs
+///
+/// This deliberatly simple interface provides the middle-end (=LLVM) with
+/// easier means to reason about the semantic of the code and transform it as
+/// well. The runtime calls are therefore also desiged to carry sufficient
+/// information necessary for optimization.
+///
+///
+/// Intended usage:
+///
+/// \code
+/// void kernel(...) {
+///
+///   char ThreadKind = __kmpc_target_region_kernel_init(...);
+///
+///   if (ThreadKind == -1) {               //  actual worker thread
+///     if (!UsedLibraryStateMachine)
+///       user_state_machine();
+///     goto exit;
+///   } else if (ThreadKind == 0) {         // surplus worker thread
+///     goto exit;
+///   } else {                              //    team master thread
+///     goto user_code;
+///   }
+///
+/// user_code:
+///
+///   // User defined kernel code, parallel regions are replaced by
+///   // by __kmpc_target_region_kernel_parallel(...) calls.
+///
+///   // Fallthrough to de-initialization
+///
+/// deinit:
+///   __kmpc_target_region_kernel_deinit(...);
+///
+/// exit:
+///   /* exit the kernel */
+/// }
+/// \endcode
+///
+///
+///{
+
+/// Initialization
+///
+/// \param Ident               Source location identification, can be NULL.
+///
+EXTERN int8_t __kmpc_device_target_init(ident_t *Ident, int32_t Num, void *);
+
+/// De-Initialization
+///
+/// In non-SPMD, this function releases the workers trapped in a state machine
+/// and also any memory dynamically allocated by the runtime.
+///
+/// \param Ident Source location identification, can be NULL.
+///
+EXTERN void __kmpc_device_target_deinit(ident_t *Ident);
+
+///}
+
+#endif
+
+EXTERN void *__kmpc_target_region_kernel_get_shared_memory() {
+  return _shared_bytes_buffer_memory.begin();
+}
+EXTERN void *__kmpc_target_region_kernel_get_private_memory() {
+  return ((char *)_shared_bytes_buffer_memory.begin()) +
+         _shared_bytes_buffer_memory.get_offset();
+}
+
+/// Simple generic state machine for worker threads.
+INLINE static void
+__kmpc_target_region_state_machine(ident_t *Ident,
+                                   bool IsOMPRuntimeInitialized) {
+
+  do {
+    void *WorkFn = 0;
+
+    // Wait for the signal that we have a new work function.
+    __kmpc_barrier_simple_spmd(Ident, 0);
+
+    // Retrieve the work function from the runtime.
+    bool IsActive = __kmpc_kernel_parallel(&WorkFn, IsOMPRuntimeInitialized);
+
+    // If there is nothing more to do, break out of the state machine by
+    // returning to the caller.
+    if (!WorkFn)
+      return;
+
+    if (IsActive) {
+      void *SharedVars = __kmpc_target_region_kernel_get_shared_memory();
+      void *PrivateVars = __kmpc_target_region_kernel_get_private_memory();
+
+      ((ParallelWorkFnTy)WorkFn)(SharedVars, PrivateVars);
+
+      __kmpc_kernel_end_parallel();
+    }
+
+    __kmpc_barrier_simple_spmd(Ident, 0);
+
+  } while (true);
+}
+
+/// Filter threads into masters and workers. If \p UseStateMachine is true,
+/// required workers will enter a state machine through and be trapped there.
+/// Master and surplus worker threads will return from this function immediately
+/// while required workers will only return once there is no more work. The
+/// return value indicates if the thread is a master (1), a surplus worker (0),
+/// or a finished required worker released from the state machine (-1).
+INLINE static int8_t
+__kmpc_target_region_thread_filter(ident_t *Ident, unsigned ThreadLimit,
+                                   bool UseStateMachine,
+                                   bool IsOMPRuntimeInitialized) {
+
+  unsigned TId = GetThreadIdInBlock();
+  bool IsWorker = TId < ThreadLimit;
+
+  if (IsWorker) {
+    if (UseStateMachine)
+      __kmpc_target_region_state_machine(Ident, IsOMPRuntimeInitialized);
+    return -1;
+  }
+
+  return TId == GetMasterThreadID();
+}
+
+EXTERN int8_t __kmpc_target_region_kernel_init(ident_t *Ident, bool UseSPMDMode,
+                                               bool RequiresOMPRuntime,
+                                               bool UseStateMachine,
+                                               bool RequiresDataSharing)
+
+    EXTERN
+    void __kmpc_target_region_kernel_deinit(ident_t *Ident, bool UseSPMDMode,
+                                            bool RequiredOMPRuntime) {
+  // Handle the SPMD case first.
+  if (UseSPMDMode) {
+    __kmpc_spmd_kernel_deinit_v2(RequiredOMPRuntime);
+    return;
+  }
+
+  __kmpc_kernel_deinit(RequiredOMPRuntime);
+
+  // Barrier to terminate worker threads.
+  __kmpc_barrier_simple_spmd(Ident, 0);
+
+  // Release any dynamically allocated memory used for sharing.
+  _shared_bytes_buffer_memory.release();
+}
+
+EXTERN void __kmpc_target_region_kernel_parallel(
+    ident_t *Ident, uint16_t UseSPMDMode, bool RequiredOMPRuntime,
+    ParallelWorkFnTy ParallelWorkFn, void *SharedVars, uint16_t SharedVarsBytes,
+    void *PrivateVars, uint16_t PrivateVarsBytes, bool SharedMemPointers) {
+
+  // If the mode is unknown we check it at runtime
+  if (UseSPMDMode == -1)
+    UseSPMDMode = __kmpc_is_spmd_exec_mode();
+
+  // In SPMD mode, we simply call the work function with the provided values.
+  if (UseSPMDMode) {
+    ParallelWorkFn(SharedVars, PrivateVars);
+    return;
+  }
+
+  if (SharedMemPointers) {
+    // If shared memory pointers are used, the user guarantees that all private
+    // variables, if any, are stored directly after the shared ones in memory.
+    // Additionally, this memory can be accessed by all the threads. In that
+    // case, we do not need to copy memory around but simply use the provided
+    // locations. However, we still need to inform the buffer of these
+    // locations as the worker threads might use the
+    //   __kmpc_target_region_kernel_get_shared_memory()
+    // and
+    //   __kmpc_target_region_kernel_get_private_memory()
+    // functions to get the respective pointers.
+
+    _shared_bytes_buffer_memory.set(SharedVars, SharedVarsBytes);
+
+  } else {
+
+    size_t BytesToCopy = SharedVarsBytes + PrivateVarsBytes;
+    if (BytesToCopy) {
+      // Resize the shared memory to be able to hold the data which is required
+      // to be in shared memory. Also set the offset to the beginning to the
+      // private variables.
+      _shared_bytes_buffer_memory.resize(BytesToCopy, SharedVarsBytes);
+
+      // Copy the shared and private variables into shared memory.
+      void *SVMemory = __kmpc_target_region_kernel_get_shared_memory();
+      void *PVMemory = __kmpc_target_region_kernel_get_private_memory();
+      memcpy(SVMemory, SharedVars, SharedVarsBytes);
+      memcpy(PVMemory, PrivateVars, PrivateVarsBytes);
+    }
+  }
+
+  // TODO: It seems we could store the work function in the same shared space
+  // as the rest of the variables above.
+  //
+  // Initialize the parallel work, e.g., make sure the work function is known.
+  __kmpc_kernel_prepare_parallel((void *)ParallelWorkFn, RequiredOMPRuntime);
+
+  // TODO: It is odd that we call the *_spmd version in non-SPMD mode here.
+  //
+  // Activate workers. This barrier is used by the master to signal
+  // work for the workers.
+  __kmpc_barrier_simple_spmd(Ident, 0);
+
+  // OpenMP [2.5, Parallel Construct, p.49]
+  // There is an implied barrier at the end of a parallel region. After the
+  // end of a parallel region, only the master thread of the team resumes
+  // execution of the enclosing task region.
+  //
+  // The master waits at this barrier until all workers are done.
+  __kmpc_barrier_simple_spmd(Ident, 0);
+
+  // Update the shared variables if necessary, that is if we did not use user
+  // memory in the first .
+  if (!SharedMemPointers && SharedVarsBytes)
+    memcpy(SharedVars, __kmpc_target_region_kernel_get_shared_memory(),
+           SharedVarsBytes);
+
+  // We could set (or reset) the _shared_bytes_buffer_memory pointer to NULL (or
+  // the old value) if we used user provided memory. This is not necessary as
+  // long as the buffer knows not to free the explicitly "set" pointer.
+}

openmp/libomptarget/deviceRTLs/common/include/target_impl.h

+//===------------ target_impl.h - 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
+//
+//===----------------------------------------------------------------------===//
+#ifndef _TARGET_IMPL_H_
+#define _TARGET_IMPL_H_
+
+#include <assert.h>
+#include <inttypes.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#define DEVICE __device__
+#define INLINE __forceinline__ DEVICE
+#define NOINLINE __noinline__ DEVICE
+#define SHARED __shared__
+
+////////////////////////////////////////////////////////////////////////////////
+// Kernel options
+////////////////////////////////////////////////////////////////////////////////
+
+void __kmpc_impl_unpack(uint64_t val, uint32_t &lo, uint32_t &hi);
+
+uint64_t __kmpc_impl_pack(uint32_t lo, uint32_t hi);
+
+// enum : __kmpc_impl_lanemask_t;
+//__kmpc_impl_all_lanes = ~(__kmpc_impl_lanemask_t)0
+//};
+
+//__kmpc_impl_lanemask_t __kmpc_impl_lanemask_lt();
+
+//__kmpc_impl_lanemask_t __kmpc_impl_lanemask_gt();
+
+uint32_t __kmpc_impl_smid();
+
+double __kmpc_impl_get_wtick();
+
+double __kmpc_impl_get_wtime();
+
+uint32_t __kmpc_impl_ffs(uint32_t x);
+
+uint32_t __kmpc_impl_popc(uint32_t x);
+
+// template <typename T> T __kmpc_impl_min(T x, T y);
+// return min(x, y);
+//}
+
+//__kmpc_impl_lanemask_t __kmpc_impl_activemask() ;
+
+// int32_t __kmpc_impl_shfl_sync(__kmpc_impl_lanemask_t Mask, int32_t Var,
+// int32_t SrcLane);
+
+// int32_t __kmpc_impl_shfl_down_sync(__kmpc_impl_lanemask_t Mask,
+// int32_t Var, uint32_t Delta,
+// int32_t Width);
+
+void __kmpc_impl_syncthreads();
+
+// void __kmpc_impl_syncwarp(__kmpc_impl_lanemask_t Mask);
+
+void __kmpc_impl_target_init();
+void __kmpc_impl_named_sync(uint32_t num_threads);
+void __kmpc_impl_threadfence(void);
+void __kmpc_impl_threadfence_block(void);
+void __kmpc_impl_threadfence_system(void);
+
+// Calls to the NVPTX layer (assuming 1D layout)
+int GetThreadIdInBlock();
+int GetBlockIdInKernel();
+int GetNumberOfBlocksInKernel();
+int GetNumberOfThreadsInBlock();
+unsigned GetWarpId();
+unsigned GetLaneId();
+
+// Locks
+struct omp_lock_t;
+void __kmpc_impl_init_lock(omp_lock_t *lock);
+void __kmpc_impl_destroy_lock(omp_lock_t *lock);
+void __kmpc_impl_set_lock(omp_lock_t *lock);
+void __kmpc_impl_unset_lock(omp_lock_t *lock);
+int __kmpc_impl_test_lock(omp_lock_t *lock);
+
+// Memory
+void *__kmpc_impl_malloc(size_t x);
+void __kmpc_impl_free(void *x);
+
+#endif

openmp/libomptarget/deviceRTLs/common/omptarget.h

 //===---- omptarget.h - OpenMP GPU initialization ---------------- 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
 //
 //===----------------------------------------------------------------------===//
 //
 // This file contains the declarations of all library macros, types,
 // and functions.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef OMPTARGET_H
 #define OMPTARGET_H
 
-#include "target_impl.h"
+//#include "target_impl.h"
 #include "common/debug.h"     // debug
 #include "interface.h" // interfaces with omp, compiler, and user
 #include "common/state-queue.h"
 #include "common/support.h"
 
+#include <stddef.h>
+#include <stdint.h>
+
 #define OMPTARGET_NVPTX_VERSION 1.1
 
 // used by the library for the interface with the app
 #define DISPATCH_FINISHED 0
 #define DISPATCH_NOTFINISHED 1
 
 // used by dynamic scheduling
 #define FINISHED 0
 #define NOT_FINISHED 1
 #define LAST_CHUNK 2
 
 #define BARRIER_COUNTER 0
 #define ORDERED_COUNTER 1
 
+#define MAX_SHARED_ARGS 20
+
 // arguments needed for L0 parallelism only.
 class omptarget_nvptx_SharedArgs {
 public:
   // All these methods must be called by the master thread only.
-  INLINE void Init() {
+  void Init() {
     args  = buffer;
     nArgs = MAX_SHARED_ARGS;
   }
-  INLINE void DeInit() {
+  void DeInit() {
     // Free any memory allocated for outlined parallel function with a large
     // number of arguments.
     if (nArgs > MAX_SHARED_ARGS) {
       SafeFree(args, "new extended args");
       Init();
     }
   }
-  INLINE void EnsureSize(size_t size) {
+  void EnsureSize(size_t size) {
     if (size > nArgs) {
       if (nArgs > MAX_SHARED_ARGS) {
         SafeFree(args, "new extended args");
       }
       args = (void **)SafeMalloc(size * sizeof(void *), "new extended args");
       nArgs = size;
     }
   }
   // Called by all threads.
-  INLINE void **GetArgs() const { return args; };
+  void **GetArgs() const { return args; };
+
 private:
   // buffer of pre-allocated arguments.
   void *buffer[MAX_SHARED_ARGS];
   // pointer to arguments buffer.
   // starts off as a pointer to 'buffer' but can be dynamically allocated.
   void **args;
   // starts off as MAX_SHARED_ARGS but can increase in size.
   uint32_t nArgs;
 };
 
-extern DEVICE SHARED omptarget_nvptx_SharedArgs
-    omptarget_nvptx_globalArgs;
+// SHARED
+extern omptarget_nvptx_SharedArgs omptarget_nvptx_globalArgs;
 
 // Worker slot type which is initialized with the default worker slot
 // size of 4*32 bytes.
 struct __kmpc_data_sharing_slot {
   __kmpc_data_sharing_slot *Next;
   __kmpc_data_sharing_slot *Prev;
   void *PrevSlotStackPtr;
   void *DataEnd;
   char Data[DS_Worker_Warp_Slot_Size];
 };
 
 // Data structure to keep in shared memory that traces the current slot, stack,
 // and frame pointer as well as the active threads that didn't exit the current
 // environment.
 struct DataSharingStateTy {
   __kmpc_data_sharing_slot *SlotPtr[DS_Max_Warp_Number];
   void *StackPtr[DS_Max_Warp_Number];
   void * volatile FramePtr[DS_Max_Warp_Number];
   __kmpc_impl_lanemask_t ActiveThreads[DS_Max_Warp_Number];
 };
 
-extern DEVICE SHARED DataSharingStateTy DataSharingState;
+// SHARED
+extern DataSharingStateTy DataSharingState;
 
 ////////////////////////////////////////////////////////////////////////////////
 // task ICV and (implicit & explicit) task state
 
 class omptarget_nvptx_TaskDescr {
 public:
   // methods for flags
-  INLINE omp_sched_t GetRuntimeSched() const;
-  INLINE void SetRuntimeSched(omp_sched_t sched);
-  INLINE int InParallelRegion() const { return items.flags & TaskDescr_InPar; }
-  INLINE int InL2OrHigherParallelRegion() const {
+  omp_sched_t GetRuntimeSched() const;
+  void SetRuntimeSched(omp_sched_t sched);
+  int InParallelRegion() const { return items.flags & TaskDescr_InPar; }
+  int InL2OrHigherParallelRegion() const {
     return items.flags & TaskDescr_InParL2P;
   }
-  INLINE int IsParallelConstruct() const {
+  int IsParallelConstruct() const {
     return items.flags & TaskDescr_IsParConstr;
   }
-  INLINE int IsTaskConstruct() const { return !IsParallelConstruct(); }
+  int IsTaskConstruct() const { return !IsParallelConstruct(); }
   // methods for other fields
-  INLINE uint16_t &ThreadId() { return items.threadId; }
-  INLINE uint64_t &RuntimeChunkSize() { return items.runtimeChunkSize; }
-  INLINE omptarget_nvptx_TaskDescr *GetPrevTaskDescr() const { return prev; }
-  INLINE void SetPrevTaskDescr(omptarget_nvptx_TaskDescr *taskDescr) {
+  uint16_t &ThreadId() { return items.threadId; }
+  uint64_t &RuntimeChunkSize() { return items.runtimeChunkSize; }
+  omptarget_nvptx_TaskDescr *GetPrevTaskDescr() const { return prev; }
+  void SetPrevTaskDescr(omptarget_nvptx_TaskDescr *taskDescr) {
     prev = taskDescr;
   }
   // init & copy
-  INLINE void InitLevelZeroTaskDescr();
-  INLINE void InitLevelOneTaskDescr(omptarget_nvptx_TaskDescr *parentTaskDescr);
-  INLINE void Copy(omptarget_nvptx_TaskDescr *sourceTaskDescr);
-  INLINE void CopyData(omptarget_nvptx_TaskDescr *sourceTaskDescr);
-  INLINE void CopyParent(omptarget_nvptx_TaskDescr *parentTaskDescr);
-  INLINE void CopyForExplicitTask(omptarget_nvptx_TaskDescr *parentTaskDescr);
-  INLINE void CopyToWorkDescr(omptarget_nvptx_TaskDescr *masterTaskDescr);
-  INLINE void CopyFromWorkDescr(omptarget_nvptx_TaskDescr *workTaskDescr);
-  INLINE void CopyConvergentParent(omptarget_nvptx_TaskDescr *parentTaskDescr,
-                                   uint16_t tid, uint16_t tnum);
-  INLINE void SaveLoopData();
-  INLINE void RestoreLoopData() const;
+  void InitLevelZeroTaskDescr();
+  void InitLevelOneTaskDescr(omptarget_nvptx_TaskDescr *parentTaskDescr);
+  void Copy(omptarget_nvptx_TaskDescr *sourceTaskDescr);
+  void CopyData(omptarget_nvptx_TaskDescr *sourceTaskDescr);
+  void CopyParent(omptarget_nvptx_TaskDescr *parentTaskDescr);
+  void CopyForExplicitTask(omptarget_nvptx_TaskDescr *parentTaskDescr);
+  void CopyToWorkDescr(omptarget_nvptx_TaskDescr *masterTaskDescr);
+  void CopyFromWorkDescr(omptarget_nvptx_TaskDescr *workTaskDescr);
+  void CopyConvergentParent(omptarget_nvptx_TaskDescr *parentTaskDescr,
+  void SaveLoopData();
+  void RestoreLoopData() const;
 
 private:
   // bits for flags: (6 used, 2 free)
   //   3 bits (SchedMask) for runtime schedule
   //   1 bit (InPar) if this thread has encountered one or more parallel region
   //   1 bit (IsParConstr) if ICV for a parallel region (false = explicit task)
   //   1 bit (InParL2+) if this thread has encountered L2 or higher parallel
   //   region

 
 ////////////////////////////////////////////////////////////////////////////////
 // Descriptor of a parallel region (worksharing in general)
 
 class omptarget_nvptx_WorkDescr {
 
 public:
   // access to data
-  INLINE omptarget_nvptx_TaskDescr *WorkTaskDescr() { return &masterTaskICV; }
+  omptarget_nvptx_TaskDescr *WorkTaskDescr() { return &masterTaskICV; }
 
 private:
   omptarget_nvptx_TaskDescr masterTaskICV;
 };
 
 ////////////////////////////////////////////////////////////////////////////////
 
 class omptarget_nvptx_TeamDescr {
 public:
   // access to data
-  INLINE omptarget_nvptx_TaskDescr *LevelZeroTaskDescr() {
+  omptarget_nvptx_TaskDescr *LevelZeroTaskDescr() {
     return &levelZeroTaskDescr;
   }
-  INLINE omptarget_nvptx_WorkDescr &WorkDescr() {
-    return workDescrForActiveParallel;
-  }
+  omptarget_nvptx_WorkDescr &WorkDescr() { return workDescrForActiveParallel; }
 
   // init
-  INLINE void InitTeamDescr();
+  void InitTeamDescr();
 
-  INLINE __kmpc_data_sharing_slot *GetPreallocatedSlotAddr(int wid) {
+  __kmpc_data_sharing_slot *GetPreallocatedSlotAddr(int wid) {
     worker_rootS[wid].DataEnd =
         &worker_rootS[wid].Data[0] + DS_Worker_Warp_Slot_Size;
     // We currently do not have a next slot.
     worker_rootS[wid].Next = 0;
     worker_rootS[wid].Prev = 0;
     worker_rootS[wid].PrevSlotStackPtr = 0;
     return (__kmpc_data_sharing_slot *)&worker_rootS[wid];
   }

 
 ////////////////////////////////////////////////////////////////////////////////
 // thread private data (struct of arrays for better coalescing)
 // tid refers here to the global thread id
 // do not support multiple concurrent kernel a this time
 class omptarget_nvptx_ThreadPrivateContext {
 public:
   // task
-  INLINE omptarget_nvptx_TaskDescr *Level1TaskDescr(int tid) {
+  omptarget_nvptx_TaskDescr *Level1TaskDescr(int tid) {
     return &levelOneTaskDescr[tid];
   }
-  INLINE void SetTopLevelTaskDescr(int tid,
-                                   omptarget_nvptx_TaskDescr *taskICV) {
+  void SetTopLevelTaskDescr(int tid, omptarget_nvptx_TaskDescr *taskICV) {
     topTaskDescr[tid] = taskICV;
   }
-  INLINE omptarget_nvptx_TaskDescr *GetTopLevelTaskDescr(int tid) const;
+  omptarget_nvptx_TaskDescr *GetTopLevelTaskDescr(int tid) const;
   // parallel
-  INLINE uint16_t &NumThreadsForNextParallel(int tid) {
-    return nextRegion.tnum[tid];
-  }
+  uint16_t &NumThreadsForNextParallel(int tid) { return nextRegion.tnum[tid]; }
   // schedule (for dispatch)
-  INLINE kmp_sched_t &ScheduleType(int tid) { return schedule[tid]; }
-  INLINE int64_t &Chunk(int tid) { return chunk[tid]; }
-  INLINE int64_t &LoopUpperBound(int tid) { return loopUpperBound[tid]; }
-  INLINE int64_t &NextLowerBound(int tid) { return nextLowerBound[tid]; }
-  INLINE int64_t &Stride(int tid) { return stride[tid]; }
+  kmp_sched_t &ScheduleType(int tid) { return schedule[tid]; }
+  int64_t &Chunk(int tid) { return chunk[tid]; }
+  int64_t &LoopUpperBound(int tid) { return loopUpperBound[tid]; }
+  int64_t &NextLowerBound(int tid) { return nextLowerBound[tid]; }
+  int64_t &Stride(int tid) { return stride[tid]; }
 
-  INLINE omptarget_nvptx_TeamDescr &TeamContext() { return teamContext; }
+  omptarget_nvptx_TeamDescr &TeamContext() { return teamContext; }
 
-  INLINE void InitThreadPrivateContext(int tid);
-  INLINE uint64_t &Cnt() { return cnt; }
+  void InitThreadPrivateContext(int tid);
+  uint64_t &Cnt() { return cnt; }
 
 private:
   // team context for this team
   omptarget_nvptx_TeamDescr teamContext;
   // task ICV for implicit threads in the only parallel region
   omptarget_nvptx_TaskDescr levelOneTaskDescr[MAX_THREADS_PER_TEAM];
   // pointer where to find the current task ICV (top of the stack)
   omptarget_nvptx_TaskDescr *topTaskDescr[MAX_THREADS_PER_TEAM];

 
 /// Memory manager for statically allocated memory.
 class omptarget_nvptx_SimpleMemoryManager {
 private:
   ALIGN(128) struct MemDataTy {
     volatile unsigned keys[OMP_STATE_COUNT];
   } MemData[MAX_SM];
 
-  INLINE static uint32_t hash(unsigned key) {
-    return key & (OMP_STATE_COUNT - 1);
-  }
+  static uint32_t hash(unsigned key) { return key & (OMP_STATE_COUNT - 1); }
 
 public:
-  INLINE void Release();
-  INLINE const void *Acquire(const void *buf, size_t size);
+  void Release();
+  const void *Acquire(const void *buf, size_t size);
 };
 
 ////////////////////////////////////////////////////////////////////////////////
 
 ////////////////////////////////////////////////////////////////////////////////
 // global data tables
 ////////////////////////////////////////////////////////////////////////////////
 
-extern DEVICE omptarget_nvptx_SimpleMemoryManager
-    omptarget_nvptx_simpleMemoryManager;
-extern DEVICE SHARED uint32_t usedMemIdx;
-extern DEVICE SHARED uint32_t usedSlotIdx;
-extern DEVICE SHARED uint8_t
-    parallelLevel[MAX_THREADS_PER_TEAM / WARPSIZE];
-extern DEVICE SHARED uint16_t threadLimit;
-extern DEVICE SHARED uint16_t threadsInTeam;
-extern DEVICE SHARED uint16_t nThreads;
-extern DEVICE SHARED
-    omptarget_nvptx_ThreadPrivateContext *omptarget_nvptx_threadPrivateContext;
+extern omptarget_nvptx_SimpleMemoryManager omptarget_nvptx_simpleMemoryManager;
+extern SHARED uint32_t usedMemIdx;
+extern SHARED uint32_t usedSlotIdx;
+extern SHARED uint8_t parallelLevel[MAX_THREADS_PER_TEAM / WARPSIZE];
+extern SHARED uint16_t threadLimit;
+extern SHARED uint16_t threadsInTeam;
+extern SHARED uint16_t nThreads;
+extern SHARED omptarget_nvptx_ThreadPrivateContext
+    *omptarget_nvptx_threadPrivateContext;
 
-extern DEVICE SHARED uint32_t execution_param;
-extern DEVICE SHARED void *ReductionScratchpadPtr;
+extern SHARED uint32_t execution_param;
+extern SHARED void *ReductionScratchpadPtr;
 
 ////////////////////////////////////////////////////////////////////////////////
 // work function (outlined parallel/simd functions) and arguments.
 // needed for L1 parallelism only.
 ////////////////////////////////////////////////////////////////////////////////
 
 typedef void *omptarget_nvptx_WorkFn;
-extern volatile DEVICE SHARED omptarget_nvptx_WorkFn
-    omptarget_nvptx_workFn;
+extern volatile SHARED omptarget_nvptx_WorkFn omptarget_nvptx_workFn;
 
 ////////////////////////////////////////////////////////////////////////////////
 // get private data structures
 ////////////////////////////////////////////////////////////////////////////////
 
-INLINE omptarget_nvptx_TeamDescr &getMyTeamDescriptor();
-INLINE omptarget_nvptx_WorkDescr &getMyWorkDescriptor();
-INLINE omptarget_nvptx_TaskDescr *
-getMyTopTaskDescriptor(bool isSPMDExecutionMode);
-INLINE omptarget_nvptx_TaskDescr *getMyTopTaskDescriptor(int globalThreadId);
+omptarget_nvptx_TeamDescr &getMyTeamDescriptor();
+omptarget_nvptx_WorkDescr &getMyWorkDescriptor();
+omptarget_nvptx_TaskDescr *getMyTopTaskDescriptor(bool isSPMDExecutionMode);
+omptarget_nvptx_TaskDescr *getMyTopTaskDescriptor(int globalThreadId);
 
 ////////////////////////////////////////////////////////////////////////////////
 // inlined implementation
 ////////////////////////////////////////////////////////////////////////////////
 
 #include "common/omptargeti.h"
 
 #endif

openmp/libomptarget/deviceRTLs/common/src/cancel.cu

 //===------ cancel.cu - NVPTX OpenMP cancel interface ------------ 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
 //
 //===----------------------------------------------------------------------===//
 //
 // Interface to be used in the implementation of OpenMP cancel.
 //
 //===----------------------------------------------------------------------===//
 
-#include "interface.h"
 #include "common/debug.h"
 
-EXTERN int32_t __kmpc_cancellationpoint(kmp_Ident *loc, int32_t global_tid,
+#include <stdint.h>
+
+struct kmp_Ident;
+
+#pragma omp declare target
+
+int32_t __kmpc_cancellationpoint(kmp_Ident *loc, int32_t global_tid,
                                         int32_t cancelVal) {
   PRINT(LD_IO, "call kmpc_cancellationpoint(cancel val %d)\n", (int)cancelVal);
   // disabled
   return 0;
 }
 
-EXTERN int32_t __kmpc_cancel(kmp_Ident *loc, int32_t global_tid,
+int32_t __kmpc_cancel(kmp_Ident *loc, int32_t global_tid,
                              int32_t cancelVal) {
   PRINT(LD_IO, "call kmpc_cancel(cancel val %d)\n", (int)cancelVal);
   // disabled
   return 0;
-}
+} 
+
+#pragma omp end declare target

openmp/libomptarget/deviceRTLs/common/src/loop.cu

 //===------------ loop.cu - NVPTX OpenMP loop constructs --------- 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
 //
 //===----------------------------------------------------------------------===//
 //
 // This file contains the implementation of the KMPC interface
 // for the loop construct plus other worksharing constructs that use the same
 // interface as loops.
 //
 //===----------------------------------------------------------------------===//
 
 #include "common/omptarget.h"
-#include "target_impl.h"
-#include "common/target_atomic.h"
+//#include "target_impl.h"
+//#include "common/target_atomic.h"
+#include "interface.h"
+#include "debug.h"
+#include "support.h"
+#include <stdint.h>
 
 ////////////////////////////////////////////////////////////////////////////////
 ////////////////////////////////////////////////////////////////////////////////
 // template class that encapsulate all the helper functions
 //
 // T is loop iteration type (32 | 64)  (unsigned | signed)
 // ST is the signed version of T
 ////////////////////////////////////////////////////////////////////////////////
public:
    *  upper bound of first chunk
    *  @param[in,out] pointer to loop stride. It will contain value of stride
    *  between two successive chunks executed by the same thread
    *  @param[in] loop increment bump
    *  @param[in] chunk size
    */
 
   // helper function for static chunk
-  INLINE static void ForStaticChunk(int &last, T &lb, T &ub, ST &stride,
+  static void ForStaticChunk(int &last, T &lb, T &ub, ST &stride,
                                     ST chunk, T entityId, T numberOfEntities) {
     // each thread executes multiple chunks all of the same size, except
     // the last one
 
     // distance between two successive chunks
     stride = numberOfEntities * chunk;
     lb = lb + entityId * chunk;
     T inputUb = ub;
     ub = lb + chunk - 1; // Clang uses i <= ub
     // Say ub' is the begining of the last chunk. Then who ever has a
     // lower bound plus a multiple of the increment equal to ub' is
     // the last one.
     T beginingLastChunk = inputUb - (inputUb % chunk);
     last = ((beginingLastChunk - lb) % stride) == 0;
   }
 
   ////////////////////////////////////////////////////////////////////////////////
   // Loop with static scheduling without chunk
 
   // helper function for static no chunk
-  INLINE static void ForStaticNoChunk(int &last, T &lb, T &ub, ST &stride,
+  static void ForStaticNoChunk(int &last, T &lb, T &ub, ST &stride,
                                       ST &chunk, T entityId,
                                       T numberOfEntities) {
     // No chunk size specified.  Each thread or warp gets at most one
     // chunk; chunks are all almost of equal size
     T loopSize = ub - lb + 1;
 
     chunk = loopSize / numberOfEntities;
     T leftOver = loopSize - chunk * numberOfEntities;
static void ForStaticNoChunk(int &last, T &lb, T &ub, ST &stride,
     ub = lb + chunk - 1; // Clang uses i <= ub
     last = lb <= inputUb && inputUb <= ub;
     stride = loopSize; // make sure we only do 1 chunk per warp
   }
 
   ////////////////////////////////////////////////////////////////////////////////
   // Support for Static Init
 
-  INLINE static void for_static_init(int32_t gtid, int32_t schedtype,
+  static void for_static_init(int32_t gtid, int32_t schedtype,
                                      int32_t *plastiter, T *plower, T *pupper,
                                      ST *pstride, ST chunk,
                                      bool IsSPMDExecutionMode) {
     // When IsRuntimeUninitialized is true, we assume that the caller is
     // in an L0 parallel region and that all worker threads participate.
 
     // Assume we are in teams region or that we use a single block
     // per target region
     ST numberOfActiveOMPThreads = GetNumberOfOmpThreads(IsSPMDExecutionMode);
PRINT(LD_LOOP,
           (int)numberOfActiveOMPThreads, (int)GetNumberOfWorkersInTeam(),
           (long long)(*plower), (long long)(*pupper), (long long)(*pstride),
           (int)lastiter);
   }
 
   ////////////////////////////////////////////////////////////////////////////////
   // Support for dispatch Init
 
-  INLINE static int OrderedSchedule(kmp_sched_t schedule) {
+  static int OrderedSchedule(kmp_sched_t schedule) {
     return schedule >= kmp_sched_ordered_first &&
            schedule <= kmp_sched_ordered_last;
   }
 
-  INLINE static void dispatch_init(kmp_Ident *loc, int32_t threadId,
+  static void dispatch_init(kmp_Ident *loc, int32_t threadId,
                                    kmp_sched_t schedule, T lb, T ub, ST st,
                                    ST chunk) {
     if (checkRuntimeUninitialized(loc)) {
       // In SPMD mode no need to check parallelism level - dynamic scheduling
       // may appear only in L2 parallel regions with lightweight runtime.
       ASSERT0(LT_FUSSY, checkSPMDMode(loc), "Expected non-SPMD mode.");
       return;
     }
     int tid = GetLogicalThreadIdInBlock(checkSPMDMode(loc));
if (schedule == kmp_sched_static_chunk) {
             omptarget_nvptx_threadPrivateContext->LoopUpperBound(tid),
             omptarget_nvptx_threadPrivateContext->Chunk(tid));
     }
   }
 
   ////////////////////////////////////////////////////////////////////////////////
   // Support for dispatch next
 
-  INLINE static uint64_t Shuffle(__kmpc_impl_lanemask_t active, int64_t val,
+  static uint64_t Shuffle(__kmpc_impl_lanemask_t active, int64_t val,
                                  int leader) {
     uint32_t lo, hi;
     __kmpc_impl_unpack(val, lo, hi);
     hi = __kmpc_impl_shfl_sync(active, hi, leader);
     lo = __kmpc_impl_shfl_sync(active, lo, leader);
     return __kmpc_impl_pack(lo, hi);
   }
 
-  INLINE static uint64_t NextIter() {
+  static uint64_t NextIter() {
     __kmpc_impl_lanemask_t active = __kmpc_impl_activemask();
     uint32_t leader = __kmpc_impl_ffs(active) - 1;
     uint32_t change = __kmpc_impl_popc(active);
     __kmpc_impl_lanemask_t lane_mask_lt = __kmpc_impl_lanemask_lt();
     unsigned int rank = __kmpc_impl_popc(active & lane_mask_lt);
     uint64_t warp_res;
     if (rank == 0) {
       warp_res = __kmpc_atomic_add(
           (unsigned long long *)&omptarget_nvptx_threadPrivateContext->Cnt(),
           (unsigned long long)change);
     }
     warp_res = Shuffle(active, warp_res, leader);
     return warp_res + rank;
   }
 
-  INLINE static int DynamicNextChunk(T &lb, T &ub, T chunkSize,
+  static int DynamicNextChunk(T &lb, T &ub, T chunkSize,
                                      T loopLowerBound, T loopUpperBound) {
     T N = NextIter();
     lb = loopLowerBound + N * chunkSize;
     ub = lb + chunkSize - 1;  // Clang uses i <= ub
 
     // 3 result cases:
     //  a. lb and ub < loopUpperBound --> NOT_FINISHED
     //  b. lb < loopUpperBound and ub >= loopUpperBound: last chunk -->
static int DynamicNextChunk(T &lb, T &ub, T chunkSize,
     // c. if we are here, we are in case 'c'
     lb = loopUpperBound + 2;
     ub = loopUpperBound + 1;
     PRINT(LD_LOOPD, "lb %lld, ub %lld, loop ub %lld; finished\n", (long long)lb,
           (long long)ub, (long long)loopUpperBound);
     return FINISHED;
   }
 
-  INLINE static int dispatch_next(kmp_Ident *loc, int32_t gtid, int32_t *plast,
+  static int dispatch_next(kmp_Ident *loc, int32_t gtid, int32_t *plast,
                                   T *plower, T *pupper, ST *pstride) {
     if (checkRuntimeUninitialized(loc)) {
       // In SPMD mode no need to check parallelism level - dynamic scheduling
       // may appear only in L2 parallel regions with lightweight runtime.
       ASSERT0(LT_FUSSY, checkSPMDMode(loc), "Expected non-SPMD mode.");
       if (*plast)
         return DISPATCH_FINISHED;
       *plast = 1;
       return DISPATCH_NOTFINISHED;
PRINT(LD_LOOP,
           "Got sched: active %d, total %d: lb %lld, ub %lld, stride = %lld, "
           "last %d\n",
           (int)GetNumberOfOmpThreads(isSPMDMode()),
           (int)GetNumberOfWorkersInTeam(), (long long)*plower,
           (long long)*pupper, (long long)*pstride, (int)*plast);
     return DISPATCH_NOTFINISHED;
   }
 
-  INLINE static void dispatch_fini() {
+  static void dispatch_fini() {
     // nothing
   }
 
   ////////////////////////////////////////////////////////////////////////////////
   // end of template class that encapsulate all the helper functions
   ////////////////////////////////////////////////////////////////////////////////
 };
 
 ////////////////////////////////////////////////////////////////////////////////
 // KMP interface implementation (dyn loops)
 ////////////////////////////////////////////////////////////////////////////////
 
 // init
-EXTERN void __kmpc_dispatch_init_4(kmp_Ident *loc, int32_t tid,
+void __kmpc_dispatch_init_4(kmp_Ident *loc, int32_t tid,
                                    int32_t schedule, int32_t lb, int32_t ub,
                                    int32_t st, int32_t chunk) {
   PRINT0(LD_IO, "call kmpc_dispatch_init_4\n");
   omptarget_nvptx_LoopSupport<int32_t, int32_t>::dispatch_init(
       loc, tid, (kmp_sched_t)schedule, lb, ub, st, chunk);
 }
 
-EXTERN void __kmpc_dispatch_init_4u(kmp_Ident *loc, int32_t tid,
+void __kmpc_dispatch_init_4u(kmp_Ident *loc, int32_t tid,
                                     int32_t schedule, uint32_t lb, uint32_t ub,
                                     int32_t st, int32_t chunk) {
   PRINT0(LD_IO, "call kmpc_dispatch_init_4u\n");
   omptarget_nvptx_LoopSupport<uint32_t, int32_t>::dispatch_init(
       loc, tid, (kmp_sched_t)schedule, lb, ub, st, chunk);
 }
 
-EXTERN void __kmpc_dispatch_init_8(kmp_Ident *loc, int32_t tid,
+void __kmpc_dispatch_init_8(kmp_Ident *loc, int32_t tid,
                                    int32_t schedule, int64_t lb, int64_t ub,
                                    int64_t st, int64_t chunk) {
   PRINT0(LD_IO, "call kmpc_dispatch_init_8\n");
   omptarget_nvptx_LoopSupport<int64_t, int64_t>::dispatch_init(
       loc, tid, (kmp_sched_t)schedule, lb, ub, st, chunk);
 }
 
-EXTERN void __kmpc_dispatch_init_8u(kmp_Ident *loc, int32_t tid,
+void __kmpc_dispatch_init_8u(kmp_Ident *loc, int32_t tid,
                                     int32_t schedule, uint64_t lb, uint64_t ub,
                                     int64_t st, int64_t chunk) {
   PRINT0(LD_IO, "call kmpc_dispatch_init_8u\n");
   omptarget_nvptx_LoopSupport<uint64_t, int64_t>::dispatch_init(
       loc, tid, (kmp_sched_t)schedule, lb, ub, st, chunk);
 }
 
 // next
-EXTERN int __kmpc_dispatch_next_4(kmp_Ident *loc, int32_t tid, int32_t *p_last,
+int __kmpc_dispatch_next_4(kmp_Ident *loc, int32_t tid, int32_t *p_last,
                                   int32_t *p_lb, int32_t *p_ub, int32_t *p_st) {
   PRINT0(LD_IO, "call kmpc_dispatch_next_4\n");
   return omptarget_nvptx_LoopSupport<int32_t, int32_t>::dispatch_next(
       loc, tid, p_last, p_lb, p_ub, p_st);
 }
 
-EXTERN int __kmpc_dispatch_next_4u(kmp_Ident *loc, int32_t tid,
+int __kmpc_dispatch_next_4u(kmp_Ident *loc, int32_t tid,
                                    int32_t *p_last, uint32_t *p_lb,
                                    uint32_t *p_ub, int32_t *p_st) {
   PRINT0(LD_IO, "call kmpc_dispatch_next_4u\n");
   return omptarget_nvptx_LoopSupport<uint32_t, int32_t>::dispatch_next(
       loc, tid, p_last, p_lb, p_ub, p_st);
 }
 
-EXTERN int __kmpc_dispatch_next_8(kmp_Ident *loc, int32_t tid, int32_t *p_last,
+int __kmpc_dispatch_next_8(kmp_Ident *loc, int32_t tid, int32_t *p_last,
                                   int64_t *p_lb, int64_t *p_ub, int64_t *p_st) {
   PRINT0(LD_IO, "call kmpc_dispatch_next_8\n");
   return omptarget_nvptx_LoopSupport<int64_t, int64_t>::dispatch_next(
       loc, tid, p_last, p_lb, p_ub, p_st);
 }
 
-EXTERN int __kmpc_dispatch_next_8u(kmp_Ident *loc, int32_t tid,
+int __kmpc_dispatch_next_8u(kmp_Ident *loc, int32_t tid,
                                    int32_t *p_last, uint64_t *p_lb,
                                    uint64_t *p_ub, int64_t *p_st) {
   PRINT0(LD_IO, "call kmpc_dispatch_next_8u\n");
   return omptarget_nvptx_LoopSupport<uint64_t, int64_t>::dispatch_next(
       loc, tid, p_last, p_lb, p_ub, p_st);
 }
 
 // fini
-EXTERN void __kmpc_dispatch_fini_4(kmp_Ident *loc, int32_t tid) {
+void __kmpc_dispatch_fini_4(kmp_Ident *loc, int32_t tid) {
   PRINT0(LD_IO, "call kmpc_dispatch_fini_4\n");
   omptarget_nvptx_LoopSupport<int32_t, int32_t>::dispatch_fini();
 }
 
-EXTERN void __kmpc_dispatch_fini_4u(kmp_Ident *loc, int32_t tid) {
+void __kmpc_dispatch_fini_4u(kmp_Ident *loc, int32_t tid) {
   PRINT0(LD_IO, "call kmpc_dispatch_fini_4u\n");
   omptarget_nvptx_LoopSupport<uint32_t, int32_t>::dispatch_fini();
 }
 
-EXTERN void __kmpc_dispatch_fini_8(kmp_Ident *loc, int32_t tid) {
+void __kmpc_dispatch_fini_8(kmp_Ident *loc, int32_t tid) {
   PRINT0(LD_IO, "call kmpc_dispatch_fini_8\n");
   omptarget_nvptx_LoopSupport<int64_t, int64_t>::dispatch_fini();
 }
 
-EXTERN void __kmpc_dispatch_fini_8u(kmp_Ident *loc, int32_t tid) {
+void __kmpc_dispatch_fini_8u(kmp_Ident *loc, int32_t tid) {
   PRINT0(LD_IO, "call kmpc_dispatch_fini_8u\n");
   omptarget_nvptx_LoopSupport<uint64_t, int64_t>::dispatch_fini();
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 // KMP interface implementation (static loops)
 ////////////////////////////////////////////////////////////////////////////////
 
-EXTERN void __kmpc_for_static_init_4(kmp_Ident *loc, int32_t global_tid,
+void __kmpc_for_static_init_4(kmp_Ident *loc, int32_t global_tid,
                                      int32_t schedtype, int32_t *plastiter,
                                      int32_t *plower, int32_t *pupper,
                                      int32_t *pstride, int32_t incr,
                                      int32_t chunk) {
   PRINT0(LD_IO, "call kmpc_for_static_init_4\n");
   omptarget_nvptx_LoopSupport<int32_t, int32_t>::for_static_init(
       global_tid, schedtype, plastiter, plower, pupper, pstride, chunk,
       checkSPMDMode(loc));
 }
 
-EXTERN void __kmpc_for_static_init_4u(kmp_Ident *loc, int32_t global_tid,
+void __kmpc_for_static_init_4u(kmp_Ident *loc, int32_t global_tid,
                                       int32_t schedtype, int32_t *plastiter,
                                       uint32_t *plower, uint32_t *pupper,
                                       int32_t *pstride, int32_t incr,
                                       int32_t chunk) {
   PRINT0(LD_IO, "call kmpc_for_static_init_4u\n");
   omptarget_nvptx_LoopSupport<uint32_t, int32_t>::for_static_init(
       global_tid, schedtype, plastiter, plower, pupper, pstride, chunk,
       checkSPMDMode(loc));
 }
 
-EXTERN void __kmpc_for_static_init_8(kmp_Ident *loc, int32_t global_tid,
+void __kmpc_for_static_init_8(kmp_Ident *loc, int32_t global_tid,
                                      int32_t schedtype, int32_t *plastiter,
                                      int64_t *plower, int64_t *pupper,
                                      int64_t *pstride, int64_t incr,
                                      int64_t chunk) {
   PRINT0(LD_IO, "call kmpc_for_static_init_8\n");
   omptarget_nvptx_LoopSupport<int64_t, int64_t>::for_static_init(
       global_tid, schedtype, plastiter, plower, pupper, pstride, chunk,
       checkSPMDMode(loc));
 }
 
-EXTERN void __kmpc_for_static_init_8u(kmp_Ident *loc, int32_t global_tid,
+void __kmpc_for_static_init_8u(kmp_Ident *loc, int32_t global_tid,
                                       int32_t schedtype, int32_t *plastiter,
                                       uint64_t *plower, uint64_t *pupper,
                                       int64_t *pstride, int64_t incr,
                                       int64_t chunk) {
   PRINT0(LD_IO, "call kmpc_for_static_init_8u\n");
   omptarget_nvptx_LoopSupport<uint64_t, int64_t>::for_static_init(
       global_tid, schedtype, plastiter, plower, pupper, pstride, chunk,
       checkSPMDMode(loc));
 }
 
-EXTERN
 void __kmpc_for_static_init_4_simple_spmd(kmp_Ident *loc, int32_t global_tid,
                                           int32_t schedtype, int32_t *plastiter,
                                           int32_t *plower, int32_t *pupper,
                                           int32_t *pstride, int32_t incr,
                                           int32_t chunk) {
   PRINT0(LD_IO, "call kmpc_for_static_init_4_simple_spmd\n");
   omptarget_nvptx_LoopSupport<int32_t, int32_t>::for_static_init(
       global_tid, schedtype, plastiter, plower, pupper, pstride, chunk,
       /*IsSPMDExecutionMode=*/true);
 }
 
-EXTERN
 void __kmpc_for_static_init_4u_simple_spmd(kmp_Ident *loc, int32_t global_tid,
                                            int32_t schedtype,
                                            int32_t *plastiter, uint32_t *plower,
                                            uint32_t *pupper, int32_t *pstride,
                                            int32_t incr, int32_t chunk) {
   PRINT0(LD_IO, "call kmpc_for_static_init_4u_simple_spmd\n");
   omptarget_nvptx_LoopSupport<uint32_t, int32_t>::for_static_init(
       global_tid, schedtype, plastiter, plower, pupper, pstride, chunk,
       /*IsSPMDExecutionMode=*/true);
 }
 
-EXTERN
 void __kmpc_for_static_init_8_simple_spmd(kmp_Ident *loc, int32_t global_tid,
                                           int32_t schedtype, int32_t *plastiter,
                                           int64_t *plower, int64_t *pupper,
                                           int64_t *pstride, int64_t incr,
                                           int64_t chunk) {
   PRINT0(LD_IO, "call kmpc_for_static_init_8_simple_spmd\n");
   omptarget_nvptx_LoopSupport<int64_t, int64_t>::for_static_init(
       global_tid, schedtype, plastiter, plower, pupper, pstride, chunk,
       /*IsSPMDExecutionMode=*/true);
 }
 
-EXTERN
 void __kmpc_for_static_init_8u_simple_spmd(kmp_Ident *loc, int32_t global_tid,
                                            int32_t schedtype,
                                            int32_t *plastiter, uint64_t *plower,
                                            uint64_t *pupper, int64_t *pstride,
                                            int64_t incr, int64_t chunk) {
   PRINT0(LD_IO, "call kmpc_for_static_init_8u_simple_spmd\n");
   omptarget_nvptx_LoopSupport<uint64_t, int64_t>::for_static_init(
       global_tid, schedtype, plastiter, plower, pupper, pstride, chunk,
       /*IsSPMDExecutionMode=*/true);
 }
 
-EXTERN
 void __kmpc_for_static_init_4_simple_generic(
     kmp_Ident *loc, int32_t global_tid, int32_t schedtype, int32_t *plastiter,
     int32_t *plower, int32_t *pupper, int32_t *pstride, int32_t incr,
     int32_t chunk) {
   PRINT0(LD_IO, "call kmpc_for_static_init_4_simple_generic\n");
   omptarget_nvptx_LoopSupport<int32_t, int32_t>::for_static_init(
       global_tid, schedtype, plastiter, plower, pupper, pstride, chunk,
       /*IsSPMDExecutionMode=*/false);
 }
 
-EXTERN
 void __kmpc_for_static_init_4u_simple_generic(
     kmp_Ident *loc, int32_t global_tid, int32_t schedtype, int32_t *plastiter,
     uint32_t *plower, uint32_t *pupper, int32_t *pstride, int32_t incr,
     int32_t chunk) {
   PRINT0(LD_IO, "call kmpc_for_static_init_4u_simple_generic\n");
   omptarget_nvptx_LoopSupport<uint32_t, int32_t>::for_static_init(
       global_tid, schedtype, plastiter, plower, pupper, pstride, chunk,
       /*IsSPMDExecutionMode=*/false);
 }
 
-EXTERN
 void __kmpc_for_static_init_8_simple_generic(
     kmp_Ident *loc, int32_t global_tid, int32_t schedtype, int32_t *plastiter,
     int64_t *plower, int64_t *pupper, int64_t *pstride, int64_t incr,
     int64_t chunk) {
   PRINT0(LD_IO, "call kmpc_for_static_init_8_simple_generic\n");
   omptarget_nvptx_LoopSupport<int64_t, int64_t>::for_static_init(
       global_tid, schedtype, plastiter, plower, pupper, pstride, chunk,
       /*IsSPMDExecutionMode=*/false);
 }
 
-EXTERN
 void __kmpc_for_static_init_8u_simple_generic(
     kmp_Ident *loc, int32_t global_tid, int32_t schedtype, int32_t *plastiter,
     uint64_t *plower, uint64_t *pupper, int64_t *pstride, int64_t incr,
     int64_t chunk) {
   PRINT0(LD_IO, "call kmpc_for_static_init_8u_simple_generic\n");
   omptarget_nvptx_LoopSupport<uint64_t, int64_t>::for_static_init(
       global_tid, schedtype, plastiter, plower, pupper, pstride, chunk,
       /*IsSPMDExecutionMode=*/false);
 }
 
-EXTERN void __kmpc_for_static_fini(kmp_Ident *loc, int32_t global_tid) {
+void __kmpc_for_static_fini(kmp_Ident *loc, int32_t global_tid) {
   PRINT0(LD_IO, "call kmpc_for_static_fini\n");
 }

openmp/libomptarget/deviceRTLs/common/state-queue.h

 //
 //===----------------------------------------------------------------------===//
 
 #ifndef __STATE_QUEUE_H
 #define __STATE_QUEUE_H
 
 #include <stdint.h>
 
-#include "target_impl.h"
+//#include "target_impl.h"
 
 template <typename ElementType, uint32_t SIZE> class omptarget_nvptx_Queue {
 private:
   ElementType elements[SIZE];
   volatile ElementType *elementQueue[SIZE];
   volatile uint32_t head;
   volatile uint32_t ids[SIZE];
   volatile uint32_t tail;
 
   static const uint32_t MAX_ID = (1u << 31) / SIZE / 2;
-  INLINE uint32_t ENQUEUE_TICKET();
-  INLINE uint32_t DEQUEUE_TICKET();
-  INLINE static uint32_t ID(uint32_t ticket);
-  INLINE bool IsServing(uint32_t slot, uint32_t id);
-  INLINE void PushElement(uint32_t slot, ElementType *element);
-  INLINE ElementType *PopElement(uint32_t slot);
-  INLINE void DoneServing(uint32_t slot, uint32_t id);
+  uint32_t ENQUEUE_TICKET();
+  uint32_t DEQUEUE_TICKET();
+  static uint32_t ID(uint32_t ticket);
+  bool IsServing(uint32_t slot, uint32_t id);
+  void PushElement(uint32_t slot, ElementType *element);
+  ElementType *PopElement(uint32_t slot);
+  void DoneServing(uint32_t slot, uint32_t id);
 
 public:
-  INLINE omptarget_nvptx_Queue() {}
-  INLINE void Enqueue(ElementType *element);
-  INLINE ElementType *Dequeue();
+  omptarget_nvptx_Queue() {}
+  void Enqueue(ElementType *element);
+  ElementType *Dequeue();
 };
 
 #include "state-queuei.h"
 
 #endif

openmp/libomptarget/deviceRTLs/common/state-queuei.h

 // Reference:
 // Thomas R.W. Scogland and Wu-chun Feng. 2015.
 // Design and Evaluation of Scalable Concurrent Queues for Many-Core
 // Architectures. International Conference on Performance Engineering.
 //
 //===----------------------------------------------------------------------===//
 
 #include "state-queue.h"
 #include "common/target_atomic.h"
 
 template <typename ElementType, uint32_t SIZE>
-INLINE uint32_t omptarget_nvptx_Queue<ElementType, SIZE>::ENQUEUE_TICKET() {
+uint32_t omptarget_nvptx_Queue<ElementType, SIZE>::ENQUEUE_TICKET() {
   return __kmpc_atomic_add((unsigned int *)&tail, 1u);
 }
 
 template <typename ElementType, uint32_t SIZE>
-INLINE uint32_t omptarget_nvptx_Queue<ElementType, SIZE>::DEQUEUE_TICKET() {
+uint32_t omptarget_nvptx_Queue<ElementType, SIZE>::DEQUEUE_TICKET() {
   return __kmpc_atomic_add((unsigned int *)&head, 1u);
 }
 
 template <typename ElementType, uint32_t SIZE>
-INLINE uint32_t
-omptarget_nvptx_Queue<ElementType, SIZE>::ID(uint32_t ticket) {
+uint32_t omptarget_nvptx_Queue<ElementType, SIZE>::ID(uint32_t ticket) {
   return (ticket / SIZE) * 2;
 }
 
 template <typename ElementType, uint32_t SIZE>
-INLINE bool omptarget_nvptx_Queue<ElementType, SIZE>::IsServing(uint32_t slot,
-                                                                uint32_t id) {
+bool omptarget_nvptx_Queue<ElementType, SIZE>::IsServing(uint32_t slot,
   return __kmpc_atomic_add((unsigned int *)&ids[slot], 0u) == id;
 }
 
 template <typename ElementType, uint32_t SIZE>
-INLINE void
-omptarget_nvptx_Queue<ElementType, SIZE>::PushElement(uint32_t slot,
-                                                      ElementType *element) {
+void omptarget_nvptx_Queue<ElementType, SIZE>::PushElement(
+    uint32_t slot, ElementType *element) {
   __kmpc_atomic_exchange((unsigned long long *)&elementQueue[slot],
                          (unsigned long long)element);
 }
 
 template <typename ElementType, uint32_t SIZE>
-INLINE ElementType *
+ElementType *
 omptarget_nvptx_Queue<ElementType, SIZE>::PopElement(uint32_t slot) {
   return (ElementType *)__kmpc_atomic_add(
       (unsigned long long *)&elementQueue[slot], (unsigned long long)0);
 }
 
 template <typename ElementType, uint32_t SIZE>
-INLINE void omptarget_nvptx_Queue<ElementType, SIZE>::DoneServing(uint32_t slot,
-                                                                  uint32_t id) {
+void omptarget_nvptx_Queue<ElementType, SIZE>::DoneServing(uint32_t slot,
   __kmpc_atomic_exchange((unsigned int *)&ids[slot], (id + 1) % MAX_ID);
 }
 
 template <typename ElementType, uint32_t SIZE>
-INLINE void
-omptarget_nvptx_Queue<ElementType, SIZE>::Enqueue(ElementType *element) {
+void omptarget_nvptx_Queue<ElementType, SIZE>::Enqueue(ElementType *element) {
   uint32_t ticket = ENQUEUE_TICKET();
   uint32_t slot = ticket % SIZE;
   uint32_t id = ID(ticket) + 1;
   while (!IsServing(slot, id))
     ;
   PushElement(slot, element);
   DoneServing(slot, id);
 }
 
 template <typename ElementType, uint32_t SIZE>
-INLINE ElementType *omptarget_nvptx_Queue<ElementType, SIZE>::Dequeue() {
+ElementType *omptarget_nvptx_Queue<ElementType, SIZE>::Dequeue() {
   uint32_t ticket = DEQUEUE_TICKET();
   uint32_t slot = ticket % SIZE;
   uint32_t id = ID(ticket);
   while (!IsServing(slot, id))
     ;
   ElementType *element = PopElement(slot);
   // This is to populate the queue because of the lack of GPU constructors.
   if (element == 0)
     element = &elements[slot];
   DoneServing(slot, id);
   return element;
 }

openmp/libomptarget/deviceRTLs/common/support.h

 //===--------- support.h - OpenMP GPU support functions ---------- 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
 //
 //===----------------------------------------------------------------------===//
 //
 // Wrapper to some functions natively supported by the GPU.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef OMPTARGET_SUPPORT_H
 #define OMPTARGET_SUPPORT_H
 
 #include "interface.h"
 #include "target_impl.h"
 
 ////////////////////////////////////////////////////////////////////////////////
 // Execution Parameters
 ////////////////////////////////////////////////////////////////////////////////
 enum ExecutionMode {
   Spmd = 0x00u,
   Generic = 0x01u,
   ModeMask = 0x01u,
 };
 
 enum RuntimeMode {
   RuntimeInitialized = 0x00u,
   RuntimeUninitialized = 0x02u,
   RuntimeMask = 0x02u,
 };
 
-DEVICE void setExecutionParameters(ExecutionMode EMode, RuntimeMode RMode);
-DEVICE bool isGenericMode();
-DEVICE bool isSPMDMode();
-DEVICE bool isRuntimeUninitialized();
-DEVICE bool isRuntimeInitialized();
+void setExecutionParameters(ExecutionMode EMode, RuntimeMode RMode);
+bool isGenericMode();
+bool isSPMDMode();
+bool isRuntimeUninitialized();
+bool isRuntimeInitialized();
 
 ////////////////////////////////////////////////////////////////////////////////
 // Execution Modes based on location parameter fields
 ////////////////////////////////////////////////////////////////////////////////
 
-DEVICE bool checkSPMDMode(kmp_Ident *loc);
-DEVICE bool checkGenericMode(kmp_Ident *loc);
-DEVICE bool checkRuntimeUninitialized(kmp_Ident *loc);
-DEVICE bool checkRuntimeInitialized(kmp_Ident *loc);
+bool checkSPMDMode(kmp_Ident *loc);
+bool checkGenericMode(kmp_Ident *loc);
+bool checkRuntimeUninitialized(kmp_Ident *loc);
+bool checkRuntimeInitialized(kmp_Ident *loc);
 
 ////////////////////////////////////////////////////////////////////////////////
 // get info from machine
 ////////////////////////////////////////////////////////////////////////////////
 
 // get global ids to locate tread/team info (constant regardless of OMP)
-DEVICE int GetLogicalThreadIdInBlock(bool isSPMDExecutionMode);
-DEVICE int GetMasterThreadID();
-DEVICE int GetNumberOfWorkersInTeam();
+int GetLogicalThreadIdInBlock(bool isSPMDExecutionMode);
+int GetMasterThreadID();
+int GetNumberOfWorkersInTeam();
 
 // get OpenMP thread and team ids
-DEVICE int GetOmpThreadId(int threadId,
-                          bool isSPMDExecutionMode);    // omp_thread_num
-DEVICE int GetOmpTeamId();                              // omp_team_num
+int GetOmpThreadId(int threadId,
+                   bool isSPMDExecutionMode); // omp_thread_num
+int GetOmpTeamId();                           // omp_team_num
 
 // get OpenMP number of threads and team
-DEVICE int GetNumberOfOmpThreads(bool isSPMDExecutionMode); // omp_num_threads
-DEVICE int GetNumberOfOmpTeams();                           // omp_num_teams
+int GetNumberOfOmpThreads(bool isSPMDExecutionMode); // omp_num_threads
+int GetNumberOfOmpTeams();                           // omp_num_teams
 
 // get OpenMP number of procs
-DEVICE int GetNumberOfProcsInTeam(bool isSPMDExecutionMode);
-DEVICE int GetNumberOfProcsInDevice(bool isSPMDExecutionMode);
+int GetNumberOfProcsInTeam(bool isSPMDExecutionMode);
+int GetNumberOfProcsInDevice(bool isSPMDExecutionMode);
 
 // masters
-DEVICE int IsTeamMaster(int ompThreadId);
+int IsTeamMaster(int ompThreadId);
 
 // Parallel level
-DEVICE void IncParallelLevel(bool ActiveParallel, __kmpc_impl_lanemask_t Mask);
-DEVICE void DecParallelLevel(bool ActiveParallel, __kmpc_impl_lanemask_t Mask);
+// void IncParallelLevel(bool ActiveParallel, __kmpc_impl_lanemask_t Mask);
+// void DecParallelLevel(bool ActiveParallel, __kmpc_impl_lanemask_t Mask);
 
 ////////////////////////////////////////////////////////////////////////////////
 // Memory
 ////////////////////////////////////////////////////////////////////////////////
 
 // safe alloc and free
-DEVICE void *SafeMalloc(size_t size, const char *msg); // check if success
-DEVICE void *SafeFree(void *ptr, const char *msg);
+void *SafeMalloc(size_t size, const char *msg); // check if success
+void *SafeFree(void *ptr, const char *msg);
 // pad to a alignment (power of 2 only)
-DEVICE unsigned long PadBytes(unsigned long size, unsigned long alignment);
+unsigned long PadBytes(unsigned long size, unsigned long alignment);
 #define ADD_BYTES(_addr, _bytes)                                               \
   ((void *)((char *)((void *)(_addr)) + (_bytes)))
 #define SUB_BYTES(_addr, _bytes)                                               \
   ((void *)((char *)((void *)(_addr)) - (_bytes)))
 
 ////////////////////////////////////////////////////////////////////////////////
 // Teams Reduction Scratchpad Helpers
 ////////////////////////////////////////////////////////////////////////////////
-DEVICE unsigned int *GetTeamsReductionTimestamp();
-DEVICE char *GetTeamsReductionScratchpad();
+unsigned int *GetTeamsReductionTimestamp();
+char *GetTeamsReductionScratchpad();
 
 #endif

openmp/libomptarget/deviceRTLs/common/target_atomic.h

 //===---- target_atomic.h - OpenMP GPU target atomic functions ---- C++ -*-===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 //
 // Declarations of atomic functions provided by each target
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef OMPTARGET_TARGET_ATOMIC_H
 #define OMPTARGET_TARGET_ATOMIC_H
 
-#include "target_impl.h"
+//#include "target_impl.h"
 
-template <typename T> INLINE T __kmpc_atomic_add(T *address, T val) {
+template <typename T> T __kmpc_atomic_add(T *address, T val) {
   return atomicAdd(address, val);
 }
 
-template <typename T> INLINE T __kmpc_atomic_inc(T *address, T val) {
+template <typename T> T __kmpc_atomic_inc(T *address, T val) {
   return atomicInc(address, val);
 }
 
-template <typename T> INLINE T __kmpc_atomic_max(T *address, T val) {
+template <typename T> T __kmpc_atomic_max(T *address, T val) {
   return atomicMax(address, val);
 }
 
-template <typename T> INLINE T __kmpc_atomic_exchange(T *address, T val) {
+template <typename T> T __kmpc_atomic_exchange(T *address, T val) {
   return atomicExch(address, val);
 }
 
-template <typename T> INLINE T __kmpc_atomic_cas(T *address, T compare, T val) {
+template <typename T> T __kmpc_atomic_cas(T *address, T compare, T val) {
   return atomicCAS(address, compare, val);
 }
 
 #endif

openmp/libomptarget/deviceRTLs/interface.h

 typedef enum omp_proc_bind_t {
   omp_proc_bind_false = 0,
   omp_proc_bind_true = 1,
   omp_proc_bind_master = 2,
   omp_proc_bind_close = 3,
   omp_proc_bind_spread = 4
 } omp_proc_bind_t;
 
-EXTERN double omp_get_wtick(void);
-EXTERN double omp_get_wtime(void);
+double omp_get_wtick(void);
+double omp_get_wtime(void);
 
-EXTERN void omp_set_num_threads(int num);
-EXTERN int omp_get_num_threads(void);
-EXTERN int omp_get_max_threads(void);
-EXTERN int omp_get_thread_limit(void);
-EXTERN int omp_get_thread_num(void);
-EXTERN int omp_get_num_procs(void);
-EXTERN int omp_in_parallel(void);
-EXTERN int omp_in_final(void);
-EXTERN void omp_set_dynamic(int flag);
-EXTERN int omp_get_dynamic(void);
-EXTERN void omp_set_nested(int flag);
-EXTERN int omp_get_nested(void);
-EXTERN void omp_set_max_active_levels(int level);
-EXTERN int omp_get_max_active_levels(void);
-EXTERN int omp_get_level(void);
-EXTERN int omp_get_active_level(void);
-EXTERN int omp_get_ancestor_thread_num(int level);
-EXTERN int omp_get_team_size(int level);
-
-EXTERN void omp_init_lock(omp_lock_t *lock);
-EXTERN void omp_init_nest_lock(omp_nest_lock_t *lock);
-EXTERN void omp_destroy_lock(omp_lock_t *lock);
-EXTERN void omp_destroy_nest_lock(omp_nest_lock_t *lock);
-EXTERN void omp_set_lock(omp_lock_t *lock);
-EXTERN void omp_set_nest_lock(omp_nest_lock_t *lock);
-EXTERN void omp_unset_lock(omp_lock_t *lock);
-EXTERN void omp_unset_nest_lock(omp_nest_lock_t *lock);
-EXTERN int omp_test_lock(omp_lock_t *lock);
-EXTERN int omp_test_nest_lock(omp_nest_lock_t *lock);
-
-EXTERN void omp_get_schedule(omp_sched_t *kind, int *modifier);
-EXTERN void omp_set_schedule(omp_sched_t kind, int modifier);
-EXTERN omp_proc_bind_t omp_get_proc_bind(void);
-EXTERN int omp_get_cancellation(void);
-EXTERN void omp_set_default_device(int deviceId);
-EXTERN int omp_get_default_device(void);
-EXTERN int omp_get_num_devices(void);
-EXTERN int omp_get_num_teams(void);
-EXTERN int omp_get_team_num(void);
-EXTERN int omp_is_initial_device(void);
-EXTERN int omp_get_initial_device(void);
-EXTERN int omp_get_max_task_priority(void);
+void omp_set_num_threads(int num);
+int omp_get_num_threads(void);
+int omp_get_max_threads(void);
+int omp_get_thread_limit(void);
+int omp_get_thread_num(void);
+int omp_get_num_procs(void);
+int omp_in_parallel(void);
+int omp_in_final(void);
+void omp_set_dynamic(int flag);
+int omp_get_dynamic(void);
+void omp_set_nested(int flag);
+int omp_get_nested(void);
+void omp_set_max_active_levels(int level);
+int omp_get_max_active_levels(void);
+int omp_get_level(void);
+int omp_get_active_level(void);
+int omp_get_ancestor_thread_num(int level);
+int omp_get_team_size(int level);
+
+struct omp_lock_t;
+void omp_init_lock(omp_lock_t *lock);
+void omp_init_nest_lock(omp_nest_lock_t *lock);
+void omp_destroy_lock(omp_lock_t *lock);
+void omp_destroy_nest_lock(omp_nest_lock_t *lock);
+void omp_set_lock(omp_lock_t *lock);
+void omp_set_nest_lock(omp_nest_lock_t *lock);
+void omp_unset_lock(omp_lock_t *lock);
+void omp_unset_nest_lock(omp_nest_lock_t *lock);
+int omp_test_lock(omp_lock_t *lock);
+int omp_test_nest_lock(omp_nest_lock_t *lock);
+
+void omp_get_schedule(omp_sched_t *kind, int *modifier);
+void omp_set_schedule(omp_sched_t kind, int modifier);
+omp_proc_bind_t omp_get_proc_bind(void);
+int omp_get_cancellation(void);
+void omp_set_default_device(int deviceId);
+int omp_get_default_device(void);
+int omp_get_num_devices(void);
+int omp_get_num_teams(void);
+int omp_get_team_num(void);
+int omp_is_initial_device(void);
+int omp_get_initial_device(void);
+int omp_get_max_task_priority(void);
 
 ////////////////////////////////////////////////////////////////////////////////
 // file below is swiped from kmpc host interface
 ////////////////////////////////////////////////////////////////////////////////
 
 ////////////////////////////////////////////////////////////////////////////////
 // kmp specific types
 ////////////////////////////////////////////////////////////////////////////////

 // sync defs
 typedef int32_t kmp_CriticalName[8];
 
 ////////////////////////////////////////////////////////////////////////////////
 // external interface
 ////////////////////////////////////////////////////////////////////////////////
 
 // parallel
-EXTERN int32_t __kmpc_global_thread_num(kmp_Ident *loc);
-EXTERN void __kmpc_push_num_threads(kmp_Ident *loc, int32_t global_tid,
-                                    int32_t num_threads);
-EXTERN void __kmpc_serialized_parallel(kmp_Ident *loc, uint32_t global_tid);
-EXTERN void __kmpc_end_serialized_parallel(kmp_Ident *loc,
-                                           uint32_t global_tid);
-EXTERN uint16_t __kmpc_parallel_level(kmp_Ident *loc, uint32_t global_tid);
+int32_t __kmpc_global_thread_num(kmp_Ident *loc);
+void __kmpc_push_num_threads(kmp_Ident *loc, int32_t global_tid,
+void __kmpc_serialized_parallel(kmp_Ident *loc, uint32_t global_tid);
+void __kmpc_end_serialized_parallel(kmp_Ident *loc, uint32_t global_tid);
+uint16_t __kmpc_parallel_level(kmp_Ident *loc, uint32_t global_tid);
 
 // proc bind
-EXTERN void __kmpc_push_proc_bind(kmp_Ident *loc, uint32_t global_tid,
-                                  int proc_bind);
-EXTERN int omp_get_num_places(void);
-EXTERN int omp_get_place_num_procs(int place_num);
-EXTERN void omp_get_place_proc_ids(int place_num, int *ids);
-EXTERN int omp_get_place_num(void);
-EXTERN int omp_get_partition_num_places(void);
-EXTERN void omp_get_partition_place_nums(int *place_nums);
+void __kmpc_push_proc_bind(kmp_Ident *loc, uint32_t global_tid, int proc_bind);
+int omp_get_num_places(void);
+int omp_get_place_num_procs(int place_num);
+void omp_get_place_proc_ids(int place_num, int *ids);
+int omp_get_place_num(void);
+int omp_get_partition_num_places(void);
+void omp_get_partition_place_nums(int *place_nums);
 
 // for static (no chunk or chunk)
-EXTERN void __kmpc_for_static_init_4(kmp_Ident *loc, int32_t global_tid,
-                                     int32_t sched, int32_t *plastiter,
-                                     int32_t *plower, int32_t *pupper,
-                                     int32_t *pstride, int32_t incr,
-                                     int32_t chunk);
-EXTERN void __kmpc_for_static_init_4u(kmp_Ident *loc, int32_t global_tid,
-                                      int32_t sched, int32_t *plastiter,
-                                      uint32_t *plower, uint32_t *pupper,
-                                      int32_t *pstride, int32_t incr,
-                                      int32_t chunk);
-EXTERN void __kmpc_for_static_init_8(kmp_Ident *loc, int32_t global_tid,
-                                     int32_t sched, int32_t *plastiter,
-                                     int64_t *plower, int64_t *pupper,
-                                     int64_t *pstride, int64_t incr,
-                                     int64_t chunk);
-EXTERN void __kmpc_for_static_init_8u(kmp_Ident *loc, int32_t global_tid,
-                                      int32_t sched, int32_t *plastiter1,
-                                      uint64_t *plower, uint64_t *pupper,
-                                      int64_t *pstride, int64_t incr,
-                                      int64_t chunk);
-EXTERN
+void __kmpc_for_static_init_4(kmp_Ident *loc, int32_t global_tid, int32_t sched,
+                              int32_t *plastiter, int32_t *plower,
+                              int32_t *pupper, int32_t *pstride, int32_t incr,
+void __kmpc_for_static_init_4u(kmp_Ident *loc, int32_t global_tid,
+                               int32_t *pstride, int32_t incr, int32_t chunk);
+void __kmpc_for_static_init_8(kmp_Ident *loc, int32_t global_tid, int32_t sched,
+                              int32_t *plastiter, int64_t *plower,
+                              int64_t *pupper, int64_t *pstride, int64_t incr,
+void __kmpc_for_static_init_8u(kmp_Ident *loc, int32_t global_tid,
+                               int64_t *pstride, int64_t incr, int64_t chunk);
 void __kmpc_for_static_init_4_simple_spmd(kmp_Ident *loc, int32_t global_tid,
                                           int32_t sched, int32_t *plastiter,
                                           int32_t *plower, int32_t *pupper,
                                           int32_t *pstride, int32_t incr,
                                           int32_t chunk);
-EXTERN
 void __kmpc_for_static_init_4u_simple_spmd(kmp_Ident *loc, int32_t global_tid,
                                            int32_t sched, int32_t *plastiter,
                                            uint32_t *plower, uint32_t *pupper,
                                            int32_t *pstride, int32_t incr,
                                            int32_t chunk);
-EXTERN
 void __kmpc_for_static_init_8_simple_spmd(kmp_Ident *loc, int32_t global_tid,
                                           int32_t sched, int32_t *plastiter,
                                           int64_t *plower, int64_t *pupper,
                                           int64_t *pstride, int64_t incr,
                                           int64_t chunk);
-EXTERN
 void __kmpc_for_static_init_8u_simple_spmd(kmp_Ident *loc, int32_t global_tid,
                                            int32_t sched, int32_t *plastiter1,
                                            uint64_t *plower, uint64_t *pupper,
                                            int64_t *pstride, int64_t incr,
                                            int64_t chunk);
-EXTERN
 void __kmpc_for_static_init_4_simple_generic(kmp_Ident *loc,
                                              int32_t global_tid, int32_t sched,
                                              int32_t *plastiter,
                                              int32_t *plower, int32_t *pupper,
                                              int32_t *pstride, int32_t incr,
                                              int32_t chunk);
-EXTERN
 void __kmpc_for_static_init_4u_simple_generic(
     kmp_Ident *loc, int32_t global_tid, int32_t sched, int32_t *plastiter,
     uint32_t *plower, uint32_t *pupper, int32_t *pstride, int32_t incr,
     int32_t chunk);
-EXTERN
 void __kmpc_for_static_init_8_simple_generic(kmp_Ident *loc,
                                              int32_t global_tid, int32_t sched,
                                              int32_t *plastiter,
                                              int64_t *plower, int64_t *pupper,
                                              int64_t *pstride, int64_t incr,
                                              int64_t chunk);
-EXTERN
 void __kmpc_for_static_init_8u_simple_generic(
     kmp_Ident *loc, int32_t global_tid, int32_t sched, int32_t *plastiter1,
     uint64_t *plower, uint64_t *pupper, int64_t *pstride, int64_t incr,
     int64_t chunk);
 
-EXTERN void __kmpc_for_static_fini(kmp_Ident *loc, int32_t global_tid);
+void __kmpc_for_static_fini(kmp_Ident *loc, int32_t global_tid);
 
 // for dynamic
-EXTERN void __kmpc_dispatch_init_4(kmp_Ident *loc, int32_t global_tid,
-                                   int32_t sched, int32_t lower, int32_t upper,
-                                   int32_t incr, int32_t chunk);
-EXTERN void __kmpc_dispatch_init_4u(kmp_Ident *loc, int32_t global_tid,
-                                    int32_t sched, uint32_t lower,
-                                    uint32_t upper, int32_t incr,
-                                    int32_t chunk);
-EXTERN void __kmpc_dispatch_init_8(kmp_Ident *loc, int32_t global_tid,
-                                   int32_t sched, int64_t lower, int64_t upper,
-                                   int64_t incr, int64_t chunk);
-EXTERN void __kmpc_dispatch_init_8u(kmp_Ident *loc, int32_t global_tid,
-                                    int32_t sched, uint64_t lower,
-                                    uint64_t upper, int64_t incr,
-                                    int64_t chunk);
+void __kmpc_dispatch_init_4(kmp_Ident *loc, int32_t global_tid, int32_t sched,
+                            int32_t lower, int32_t upper, int32_t incr,
+                            int32_t chunk);
+void __kmpc_dispatch_init_4u(kmp_Ident *loc, int32_t global_tid, int32_t sched,
+                             uint32_t lower, uint32_t upper, int32_t incr,
+void __kmpc_dispatch_init_8(kmp_Ident *loc, int32_t global_tid, int32_t sched,
+                            int64_t lower, int64_t upper, int64_t incr,
+                            int64_t chunk);
+void __kmpc_dispatch_init_8u(kmp_Ident *loc, int32_t global_tid, int32_t sched,
+                             uint64_t lower, uint64_t upper, int64_t incr,
 
-EXTERN int __kmpc_dispatch_next_4(kmp_Ident *loc, int32_t global_tid,
-                                  int32_t *plastiter, int32_t *plower,
-                                  int32_t *pupper, int32_t *pstride);
-EXTERN int __kmpc_dispatch_next_4u(kmp_Ident *loc, int32_t global_tid,
-                                   int32_t *plastiter, uint32_t *plower,
-                                   uint32_t *pupper, int32_t *pstride);
-EXTERN int __kmpc_dispatch_next_8(kmp_Ident *loc, int32_t global_tid,
-                                  int32_t *plastiter, int64_t *plower,
-                                  int64_t *pupper, int64_t *pstride);
-EXTERN int __kmpc_dispatch_next_8u(kmp_Ident *loc, int32_t global_tid,
-                                   int32_t *plastiter, uint64_t *plower,
-                                   uint64_t *pupper, int64_t *pstride);
+int __kmpc_dispatch_next_4(kmp_Ident *loc, int32_t global_tid,
+                           int32_t *plastiter, int32_t *plower, int32_t *pupper,
+                           int32_t *pstride);
+int __kmpc_dispatch_next_4u(kmp_Ident *loc, int32_t global_tid,
+int __kmpc_dispatch_next_8(kmp_Ident *loc, int32_t global_tid,
+                           int32_t *plastiter, int64_t *plower, int64_t *pupper,
+                           int64_t *pstride);
+int __kmpc_dispatch_next_8u(kmp_Ident *loc, int32_t global_tid,
 
-EXTERN void __kmpc_dispatch_fini_4(kmp_Ident *loc, int32_t global_tid);
-EXTERN void __kmpc_dispatch_fini_4u(kmp_Ident *loc, int32_t global_tid);
-EXTERN void __kmpc_dispatch_fini_8(kmp_Ident *loc, int32_t global_tid);
-EXTERN void __kmpc_dispatch_fini_8u(kmp_Ident *loc, int32_t global_tid);
+void __kmpc_dispatch_fini_4(kmp_Ident *loc, int32_t global_tid);
+void __kmpc_dispatch_fini_4u(kmp_Ident *loc, int32_t global_tid);
+void __kmpc_dispatch_fini_8(kmp_Ident *loc, int32_t global_tid);
+void __kmpc_dispatch_fini_8u(kmp_Ident *loc, int32_t global_tid);
 
 // reduction
-EXTERN void __kmpc_nvptx_end_reduce(int32_t global_tid);
-EXTERN void __kmpc_nvptx_end_reduce_nowait(int32_t global_tid);
-EXTERN int32_t __kmpc_nvptx_parallel_reduce_nowait_v2(
+void __kmpc_nvptx_end_reduce(int32_t global_tid);
+void __kmpc_nvptx_end_reduce_nowait(int32_t global_tid);
+int32_t __kmpc_nvptx_parallel_reduce_nowait_v2(
     kmp_Ident *loc, int32_t global_tid, int32_t num_vars, size_t reduce_size,
     void *reduce_data, kmp_ShuffleReductFctPtr shflFct,
     kmp_InterWarpCopyFctPtr cpyFct);
-EXTERN int32_t __kmpc_nvptx_teams_reduce_nowait_v2(
+int32_t __kmpc_nvptx_teams_reduce_nowait_v2(
     kmp_Ident *loc, int32_t global_tid, void *global_buffer,
     int32_t num_of_records, void *reduce_data, kmp_ShuffleReductFctPtr shflFct,
     kmp_InterWarpCopyFctPtr cpyFct, kmp_ListGlobalFctPtr lgcpyFct,
     kmp_ListGlobalFctPtr lgredFct, kmp_ListGlobalFctPtr glcpyFct,
     kmp_ListGlobalFctPtr glredFct);
-EXTERN int32_t __kmpc_shuffle_int32(int32_t val, int16_t delta, int16_t size);
-EXTERN int64_t __kmpc_shuffle_int64(int64_t val, int16_t delta, int16_t size);
+int32_t __kmpc_shuffle_int32(int32_t val, int16_t delta, int16_t size);
+int64_t __kmpc_shuffle_int64(int64_t val, int16_t delta, int16_t size);
 
 // sync barrier
-EXTERN void __kmpc_barrier(kmp_Ident *loc_ref, int32_t tid);
-EXTERN void __kmpc_barrier_simple_spmd(kmp_Ident *loc_ref, int32_t tid);
-EXTERN int32_t __kmpc_cancel_barrier(kmp_Ident *loc, int32_t global_tid);
+void __kmpc_barrier(kmp_Ident *loc_ref, int32_t tid);
+void __kmpc_barrier_simple_spmd(kmp_Ident *loc_ref, int32_t tid);
+int32_t __kmpc_cancel_barrier(kmp_Ident *loc, int32_t global_tid);
 
 // single
-EXTERN int32_t __kmpc_single(kmp_Ident *loc, int32_t global_tid);
-EXTERN void __kmpc_end_single(kmp_Ident *loc, int32_t global_tid);
+int32_t __kmpc_single(kmp_Ident *loc, int32_t global_tid);
+void __kmpc_end_single(kmp_Ident *loc, int32_t global_tid);
 
 // sync
-EXTERN int32_t __kmpc_master(kmp_Ident *loc, int32_t global_tid);
-EXTERN void __kmpc_end_master(kmp_Ident *loc, int32_t global_tid);
-EXTERN void __kmpc_ordered(kmp_Ident *loc, int32_t global_tid);
-EXTERN void __kmpc_end_ordered(kmp_Ident *loc, int32_t global_tid);
-EXTERN void __kmpc_critical(kmp_Ident *loc, int32_t global_tid,
-                            kmp_CriticalName *crit);
-EXTERN void __kmpc_end_critical(kmp_Ident *loc, int32_t global_tid,
-                                kmp_CriticalName *crit);
-EXTERN void __kmpc_flush(kmp_Ident *loc);
+int32_t __kmpc_master(kmp_Ident *loc, int32_t global_tid);
+void __kmpc_end_master(kmp_Ident *loc, int32_t global_tid);
+void __kmpc_ordered(kmp_Ident *loc, int32_t global_tid);
+void __kmpc_end_ordered(kmp_Ident *loc, int32_t global_tid);
+void __kmpc_critical(kmp_Ident *loc, int32_t global_tid,
+void __kmpc_end_critical(kmp_Ident *loc, int32_t global_tid,
+void __kmpc_flush(kmp_Ident *loc);
 
 // vote
-EXTERN __kmpc_impl_lanemask_t __kmpc_warp_active_thread_mask();
+//__kmpc_impl_lanemask_t __kmpc_warp_active_thread_mask();
 // syncwarp
-EXTERN void __kmpc_syncwarp(__kmpc_impl_lanemask_t);
+// void __kmpc_syncwarp(__kmpc_impl_lanemask_t);
 
 // tasks
-EXTERN kmp_TaskDescr *__kmpc_omp_task_alloc(kmp_Ident *loc,
-                                            uint32_t global_tid, int32_t flag,
-                                            size_t sizeOfTaskInclPrivate,
-                                            size_t sizeOfSharedTable,
-                                            kmp_TaskFctPtr sub);
-EXTERN int32_t __kmpc_omp_task(kmp_Ident *loc, uint32_t global_tid,
-                               kmp_TaskDescr *newLegacyTaskDescr);
-EXTERN int32_t __kmpc_omp_task_with_deps(kmp_Ident *loc, uint32_t global_tid,
-                                         kmp_TaskDescr *newLegacyTaskDescr,
-                                         int32_t depNum, void *depList,
-                                         int32_t noAliasDepNum,
-                                         void *noAliasDepList);
-EXTERN void __kmpc_omp_task_begin_if0(kmp_Ident *loc, uint32_t global_tid,
-                                      kmp_TaskDescr *newLegacyTaskDescr);
-EXTERN void __kmpc_omp_task_complete_if0(kmp_Ident *loc, uint32_t global_tid,
-                                         kmp_TaskDescr *newLegacyTaskDescr);
-EXTERN void __kmpc_omp_wait_deps(kmp_Ident *loc, uint32_t global_tid,
-                                 int32_t depNum, void *depList,
-                                 int32_t noAliasDepNum, void *noAliasDepList);
-EXTERN void __kmpc_taskgroup(kmp_Ident *loc, uint32_t global_tid);
-EXTERN void __kmpc_end_taskgroup(kmp_Ident *loc, uint32_t global_tid);
-EXTERN int32_t __kmpc_omp_taskyield(kmp_Ident *loc, uint32_t global_tid,
-                                    int end_part);
-EXTERN int32_t __kmpc_omp_taskwait(kmp_Ident *loc, uint32_t global_tid);
-EXTERN void __kmpc_taskloop(kmp_Ident *loc, uint32_t global_tid,
-                            kmp_TaskDescr *newKmpTaskDescr, int if_val,
-                            uint64_t *lb, uint64_t *ub, int64_t st, int nogroup,
-                            int32_t sched, uint64_t grainsize, void *task_dup);
+kmp_TaskDescr *__kmpc_omp_task_alloc(kmp_Ident *loc, uint32_t global_tid,
+                                     int32_t flag, size_t sizeOfTaskInclPrivate,
+int32_t __kmpc_omp_task(kmp_Ident *loc, uint32_t global_tid,
+int32_t __kmpc_omp_task_with_deps(kmp_Ident *loc, uint32_t global_tid,
+                                  int32_t noAliasDepNum, void *noAliasDepList);
+void __kmpc_omp_task_begin_if0(kmp_Ident *loc, uint32_t global_tid,
+void __kmpc_omp_task_complete_if0(kmp_Ident *loc, uint32_t global_tid,
+void __kmpc_omp_wait_deps(kmp_Ident *loc, uint32_t global_tid, int32_t depNum,
+                          void *depList, int32_t noAliasDepNum,
+                          void *noAliasDepList);
+void __kmpc_taskgroup(kmp_Ident *loc, uint32_t global_tid);
+void __kmpc_end_taskgroup(kmp_Ident *loc, uint32_t global_tid);
+int32_t __kmpc_omp_taskyield(kmp_Ident *loc, uint32_t global_tid, int end_part);
+int32_t __kmpc_omp_taskwait(kmp_Ident *loc, uint32_t global_tid);
+void __kmpc_taskloop(kmp_Ident *loc, uint32_t global_tid,
+                     kmp_TaskDescr *newKmpTaskDescr, int if_val, uint64_t *lb,
+                     uint64_t *ub, int64_t st, int nogroup, int32_t sched,
+                     uint64_t grainsize, void *task_dup);
 
 // cancel
-EXTERN int32_t __kmpc_cancellationpoint(kmp_Ident *loc, int32_t global_tid,
-                                        int32_t cancelVal);
-EXTERN int32_t __kmpc_cancel(kmp_Ident *loc, int32_t global_tid,
+int32_t __kmpc_cancellationpoint(kmp_Ident *loc, int32_t global_tid,
                                  int32_t cancelVal);
+int32_t __kmpc_cancel(kmp_Ident *loc, int32_t global_tid, int32_t cancelVal);
 
 // non standard
-EXTERN void __kmpc_kernel_init(int ThreadLimit, int16_t RequiresOMPRuntime);
-EXTERN void __kmpc_kernel_deinit(int16_t IsOMPRuntimeInitialized);
-EXTERN void __kmpc_spmd_kernel_init(int ThreadLimit,
-                                    int16_t RequiresOMPRuntime);
-EXTERN void __kmpc_spmd_kernel_deinit_v2(int16_t RequiresOMPRuntime);
-EXTERN void __kmpc_kernel_prepare_parallel(void *WorkFn);
-EXTERN bool __kmpc_kernel_parallel(void **WorkFn);
-EXTERN void __kmpc_kernel_end_parallel();
-
-EXTERN void __kmpc_data_sharing_init_stack();
-EXTERN void __kmpc_data_sharing_init_stack_spmd();
-EXTERN void *__kmpc_data_sharing_coalesced_push_stack(size_t size,
+void __kmpc_kernel_init(int ThreadLimit, int16_t RequiresOMPRuntime);
+void __kmpc_kernel_deinit(int16_t IsOMPRuntimeInitialized);
+void __kmpc_spmd_kernel_init(int ThreadLimit, int16_t RequiresOMPRuntime);
+void __kmpc_spmd_kernel_deinit_v2(int16_t RequiresOMPRuntime);
+void __kmpc_kernel_prepare_parallel(void *WorkFn);
+bool __kmpc_kernel_parallel(void **WorkFn);
+void __kmpc_kernel_end_parallel();
+
+void __kmpc_data_sharing_init_stack();
+void __kmpc_data_sharing_init_stack_spmd();
+void *__kmpc_data_sharing_coalesced_push_stack(size_t size,
                                                int16_t UseSharedMemory);
-EXTERN void *__kmpc_data_sharing_push_stack(size_t size, int16_t UseSharedMemory);
-EXTERN void __kmpc_data_sharing_pop_stack(void *a);
-EXTERN void __kmpc_begin_sharing_variables(void ***GlobalArgs, size_t nArgs);
-EXTERN void __kmpc_end_sharing_variables();
-EXTERN void __kmpc_get_shared_variables(void ***GlobalArgs);
+void *__kmpc_data_sharing_push_stack(size_t size, int16_t UseSharedMemory);
+void __kmpc_data_sharing_pop_stack(void *a);
+void __kmpc_begin_sharing_variables(void ***GlobalArgs, size_t nArgs);
+void __kmpc_end_sharing_variables();
+void __kmpc_get_shared_variables(void ***GlobalArgs);
 
 // SPMD execution mode interrogation function.
-EXTERN int8_t __kmpc_is_spmd_exec_mode();
+int8_t __kmpc_is_spmd_exec_mode();
 
-EXTERN void __kmpc_get_team_static_memory(int16_t isSPMDExecutionMode,
-                                          const void *buf, size_t size,
-                                          int16_t is_shared, const void **res);
+void __kmpc_get_team_static_memory(int16_t isSPMDExecutionMode, const void *buf,
+                                   size_t size, int16_t is_shared,
+                                   const void **res);
 
-EXTERN void __kmpc_restore_team_static_memory(int16_t isSPMDExecutionMode,
-                                              int16_t is_shared);
+void __kmpc_restore_team_static_memory(int16_t isSPMDExecutionMode,
 
 #endif