[OpenMP][PoC][WIP] An OpenMP-based OpenMP device runtime [NOT FOR COMMIT]

[jdoerfert]

Summary

In order to make the device runtime truly device independent and thereby
reusable, e.g., for testing on the host, we want to eliminate the
remaining non-portable code. This basically boils down to device,
shared, and some other attributes, e.g., forceinline.

NOTE: This does not yet create a valid .bc file as the IR files we get via c++ are not directly usable by llvm-link/opt/.... I did not investigate how to extract the pure IR. The entire nvptx compilation process should be re-investigated as we might not need to link anything from cuda in.

[jdoerfert]

Use pragma syntax for allocators (e.g., to replace shared)

[jdoerfert]

We moved to a CUDA-free runtime now.

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 "common/target.h"
 
 struct omptarget_device_environmentTy {
   int32_t debug_level;
 };
 
-extern DEVICE omptarget_device_environmentTy omptarget_device_environment;
+__LEAGUE_VAR(extern omptarget_device_environmentTy, omptarget_device_environment)
 
 #endif

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

+//===------ cancel.cpp - NVPTX OpenMP cancel interface ------------ 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
+//
+//===----------------------------------------------------------------------===//
+//
+// Interface to be used in the implementation of OpenMP cancel.
+//
+//===----------------------------------------------------------------------===//
+
+//#include "common/debug.h"
+#include "common/target.h"
+
+struct kmp_Ident;
+
+__DEVICE_SCOPE_BEGIN()
+
+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 1;
+}
+
+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;
+}
+
+__DEVICE_SCOPE_END()

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,
-                                        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 cancelVal) {
-  PRINT(LD_IO, "call kmpc_cancel(cancel val %d)\n", (int)cancelVal);
-  // disabled
-  return 0;
-}

openmp/libomptarget/deviceRTLs/common/src/critical.cpp

-//===------ critical.cu - NVPTX OpenMP critical ------------------ CUDA -*-===//
+//===------ critical.cpp - NVPTX OpenMP critical ------------------ 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
 //
 //===----------------------------------------------------------------------===//
 //
 // This file contains the implementation of critical with KMPC interface
 //
 //===----------------------------------------------------------------------===//
 
+//#include "common/debug.h"
+#include "nvptx_interface.h"
 #include "interface.h"
-#include "common/debug.h"
 
-EXTERN
+#include "common/target.h"
+
+__DEVICE_SCOPE_BEGIN()
+
 void __kmpc_critical(kmp_Ident *loc, int32_t global_tid,
                      kmp_CriticalName *lck) {
-  PRINT0(LD_IO, "call to kmpc_critical()\n");
+  //PRINT0(LD_IO, "call to kmpc_critical()\n");
   omp_set_lock((omp_lock_t *)lck);
 }
 
-EXTERN
 void __kmpc_end_critical(kmp_Ident *loc, int32_t global_tid,
                          kmp_CriticalName *lck) {
-  PRINT0(LD_IO, "call to kmpc_end_critical()\n");
+  //PRINT0(LD_IO, "call to kmpc_end_critical()\n");
   omp_unset_lock((omp_lock_t *)lck);
 }
+
+__DEVICE_SCOPE_END()

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"
+#include "interface.h"
 
 ////////////////////////////////////////////////////////////////////////////////
 // Execution Parameters
 ////////////////////////////////////////////////////////////////////////////////
 enum ExecutionMode {
   Spmd = 0x00u,
   Generic = 0x01u,
   ModeMask = 0x01u,

openmp/libomptarget/deviceRTLs/common/target.h

+//===---- target.h - OpenMP defines and helpers for target code --- 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
+//
+//===----------------------------------------------------------------------===//
+//
+// Defines and helpers for target code.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef OMPTARGET_TARGET_H
+#define OMPTARGET_TARGET_H
+
+#include <cstdint>
+
+#define __p(STR) _Pragma(STR)
+#define __p2(STR) __p(#STR)
+
+#define __DEVICE_SCOPE_BEGIN()                                                 \
+  extern "C" {                                                                 \
+  __p("omp declare target")
+
+#define __DEVICE_SCOPE_END()                                                   \
+  __p("omp end declare target")                                                \
+  } /* extern "C" */
+
+#define __CONSTEXPR static constexpr __attribute__((nothrow, always_inline))
+
+#define __LEAGUE_VAR(TYPE, NAME)                                               \
+  TYPE NAME [[clang::loader_uninitialized]];                                   \
+  __p2(omp declare target to(NAME))
+
+#define __TEAM_VAR(TYPE, NAME)                                                 \
+  TYPE NAME [[clang::loader_uninitialized]];                                   \
+  __p2(omp allocate(NAME) allocator(omp_pteam_mem_alloc))                      \
+      __p2(omp declare target to(NAME))
+
+#define __THREAD_VAR(TYPE, NAME)                                               \
+  TYPE NAME [[clang::loader_uninitialized]];                                   \
+  __p2(omp allocate(NAME) allocator(omp_thread_mem_alloc))                     \
+      __p2(omp declare target to(NAME))
+
+#endif

openmp/libomptarget/deviceRTLs/nvptx/CMakeLists.txt

 
 if(LIBOMPTARGET_DEP_CUDA_FOUND)
   libomptarget_say("Building CUDA offloading device RTL.")
 
   # We really don't have any host code, so we don't need to care about
   # propagating host flags.
   set(CUDA_PROPAGATE_HOST_FLAGS OFF)
 
+  set(cxx_src_files
+      ${devicertl_common_directory}/src/cancel.cpp
+      ${devicertl_common_directory}/src/critical.cpp
+  )
+
   set(cuda_src_files
-      ${devicertl_common_directory}/src/cancel.cu
-      ${devicertl_common_directory}/src/critical.cu
       ${devicertl_common_directory}/src/data_sharing.cu
       ${devicertl_common_directory}/src/libcall.cu
       ${devicertl_common_directory}/src/loop.cu
       ${devicertl_common_directory}/src/omptarget.cu
       ${devicertl_common_directory}/src/parallel.cu
       ${devicertl_common_directory}/src/reduction.cu
       ${devicertl_common_directory}/src/support.cu
       ${devicertl_common_directory}/src/sync.cu
if(LIBOMPTARGET_DEP_CUDA_FOUND)
   if(${LIBOMPTARGET_NVPTX_DEBUG})
     set(CUDA_DEBUG -DOMPTARGET_NVPTX_DEBUG=-1 -g --ptxas-options=-v)
   endif()
 
   # NVPTX runtime library has to be statically linked. Dynamic linking is not
   # yet supported by the CUDA toolchain on the device.
   set(BUILD_SHARED_LIBS OFF)
   set(CUDA_SEPARABLE_COMPILATION ON)
+  list(APPEND CXX_FLAGS -I${devicertl_base_directory}
+                        -I${devicertl_nvptx_directory}/src)
   list(APPEND CUDA_NVCC_FLAGS -I${devicertl_base_directory}
                               -I${devicertl_nvptx_directory}/src)
   cuda_add_library(omptarget-nvptx STATIC ${cuda_src_files} ${omp_data_objects}
       OPTIONS ${CUDA_ARCH} ${CUDA_DEBUG})
 
   # Install device RTL under the lib destination folder.
   install(TARGETS omptarget-nvptx ARCHIVE DESTINATION "${OPENMP_INSTALL_LIBDIR}")
 
if (${LIBOMPTARGET_NVPTX_ENABLE_BCLIB})
 
     # Create target to build all Bitcode libraries.
     add_custom_target(omptarget-nvptx-bc)
 
     # Generate a Bitcode library for all the compute capabilities the user requested.
     foreach(sm ${nvptx_sm_list})
       set(cuda_arch --cuda-gpu-arch=sm_${sm})
 
-      # Compile CUDA files to bitcode.
+      # Compile C++ files to bitcode.
       set(bc_files "")
+      foreach(src ${cxx_src_files})
+        get_filename_component(infile ${src} ABSOLUTE)
+        get_filename_component(outfile ${src} NAME)
+
+        add_custom_command(OUTPUT ${outfile}-sm_${sm}.bc
+          COMMAND ${LIBOMPTARGET_NVPTX_SELECTED_CUDA_COMPILER} ${bc_flags} ${cuda_arch}
+            -c ${infile} -o ${outfile}-sm_${sm}.bc -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda
+            -Xclang -fopenmp-is-device
+            -Xclang -aux-triple -Xclang x86_64-unknown-linux
+          DEPENDS ${infile}
+          IMPLICIT_DEPENDS CXX ${infile}
+          COMMENT "Building LLVM bitcode ${outfile}-sm_${sm}.bc"
+          VERBATIM
+        )
+        set_property(DIRECTORY APPEND PROPERTY ADDITIONAL_MAKE_CLEAN_FILES ${outfile}-sm_${sm}.bc)
+
+        list(APPEND bc_files ${outfile}-sm_${sm}.bc)
+      endforeach()
+
+      # Compile CUDA files to bitcode.
       foreach(src ${cuda_src_files})
         get_filename_component(infile ${src} ABSOLUTE)
         get_filename_component(outfile ${src} NAME)
 
         add_custom_command(OUTPUT ${outfile}-sm_${sm}.bc
           COMMAND ${LIBOMPTARGET_NVPTX_SELECTED_CUDA_COMPILER} ${bc_flags} ${cuda_arch}
             -c ${infile} -o ${outfile}-sm_${sm}.bc
           DEPENDS ${infile}

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

 //===--- nvptx_interface.h - OpenMP interface definitions -------- CUDA -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef _NVPTX_INTERFACE_H_
 #define _NVPTX_INTERFACE_H_
 
 #include <stdint.h>
 
+#ifndef _OPENMP
 #define EXTERN extern "C" __device__
+#else
+#define EXTERN extern "C"
+#endif
+
 typedef uint32_t __kmpc_impl_lanemask_t;
 typedef uint32_t omp_lock_t; /* arbitrary type of the right length */
 
 #endif