[libomptarget][nfc] Make interface.h target independent

[JonChesterfield]

Summary

[libomptarget][nfc] Make interface.h target independent

Move interface.h under a top level include directory.
Remove #includes to avoid the interface depending on the implementation.

[JonChesterfield]

openmp/libomptarget/deviceRTLs/interface.h

20

__CUDACC__ seems the preferred way of spelling this constraint. The amdgcn instantiation would add an equivalent few lines to this header and otherwise leave it unchanged:

#ifdef __AMDGCN__
#define EXTERN extern "C" __attribute__((device))
typedef uint64_t __kmpc_impl_lanemask_t;
#endif

openmp/libomptarget/deviceRTLs/nvptx/CMakeLists.txt

37–41

Google suggests this is idiomatic cmake for 'find parent directory'

[jdoerfert]

I like that we start to extract common code out of .../nvptx but I'm confused about the how. As I mention below, I expected "all generic code" to include target_impl.h in order to avoid target specific code (here the __device__ stuff) under guards. Maybe I miss something so please let me know what you think.

---

I was expecting something like:
include/{interface.h, option.h, ...} all contain an #include "target_impl.h" at the beginning (see below on two options for that).
Code like the __kmpc_impl_lanemask_t typedef and defines like NOINLINE are in this setup part of target_impl.h.

openmp/libomptarget/deviceRTLs/interface.h

24

I would have thought we include target_impl.h here to get the target specific stuff.
Either via

#ifdef __CUDACC__
#include "nvptx/target_impl.h"
#endif

or via cmake include path magic.

[JonChesterfield]

I'm expecting most code moved under common to need #include "target_impl.h", located as you describe. loop.cu etc fit that category, a couple of files may not need any target stuff at all.

The library API is a flat description of everything the user (well, mostly the compiler) can find within the compiled library. It's currently called interface.h (I'd prefer libomptarget/deviceRTLs/deviceRTL.h). Including options.h or target_impl.h in the interface (as we have now) pulls in a bunch of macros and symbols that happen to be nearby.

Equivalent to this change would be creating a file nvptx/src/target_interface.h (or maybe nvptx_interface.h) containing the #define and typedef, then changing interface.h to:

#ifdef __CUDACC__
#include "../nvptx/src/target_interface.h"
#endif

Or s$src$include$, and optionally replace the relative include with more cmake.

I'd like to be able to #include "deviceRTL/interface.h" from test code and have no idea which instantiation it runs against, but it's not a disaster if the interface.h also needs a bunch of cmake wiring to locate the rest of said interface. My main goal here is to get the implementation symbols out of the library interface.

[JonChesterfield]

• Move interface details into separate file

[JonChesterfield]

Sketch two. Thanks for the feedback - I like the separate file more. It provides an obvious place to put architecture specific extensions to the API. Avoided relative includes by dropping the include dir.

[jdoerfert]

I like this better. I'll review it tomorrow (I hope).

[JonChesterfield]

Ping. I'm hoping this is uncontentious.

[jdoerfert]

In D68615#1708377, @JonChesterfield wrote:

Ping. I'm hoping this is uncontentious.

LGTM.

openmp/libomptarget/deviceRTLs/interface.h

-//===------- interface.h - NVPTX OpenMP interface definitions ---- CUDA -*-===//
+//===------- 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
 //
 //===----------------------------------------------------------------------===//
 //
-// This file contains debug macros to be used in the application.
-//
 //  This file contains all the definitions that are relevant to
 //  the interface. The first section contains the interface as
 //  declared by OpenMP.  The second section includes the compiler
 //  specific interfaces.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef _INTERFACES_H_
 #define _INTERFACES_H_
 
-#include "option.h"
-#include "target_impl.h"
+#include <stdint.h>
+
+#ifdef __CUDACC__
+#include "nvptx/src/nvptx_interface.h"
+#endif
 
 ////////////////////////////////////////////////////////////////////////////////
 // OpenMP interface
 ////////////////////////////////////////////////////////////////////////////////
 
 typedef uint32_t omp_lock_t;      /* arbitrary type of the right length */
 typedef uint64_t omp_nest_lock_t; /* arbitrary type of the right length */
 
 typedef enum omp_sched_t {

openmp/libomptarget/deviceRTLs/nvptx/CMakeLists.txt

if(CUDA_HOST_COMPILER MATCHES clang)
 
   if(NOT LIBOMPTARGET_NVPTX_ALTERNATE_GCC_HOST_COMPILER)
     libomptarget_say("Not building CUDA offloading device RTL: clang is not supported as NVCC host compiler.")
     libomptarget_say("Please include gcc in your path or set LIBOMPTARGET_NVPTX_ALTERNATE_HOST_COMPILER to the full path of of valid compiler.")
     return()
   endif()
   set(CUDA_HOST_COMPILER "${LIBOMPTARGET_NVPTX_ALTERNATE_GCC_HOST_COMPILER}" CACHE FILEPATH "" FORCE)
 endif()
 
+get_filename_component(devicertl_base_directory
+  ${CMAKE_CURRENT_SOURCE_DIR}
+  DIRECTORY)
+
 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(cuda_src_files
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 CUDA_NVCC_FLAGS -I${devicertl_base_directory})
   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}")
 
   target_link_libraries(omptarget-nvptx ${CUDA_LIBRARIES})
 
set(LIBOMPTARGET_NVPTX_ENABLE_BCLIB ${bclib_default} CACHE BOOL
     "Enable CUDA LLVM bitcode offloading device RTL.")
   if (${LIBOMPTARGET_NVPTX_ENABLE_BCLIB})
     if (NOT ${LIBOMPTARGET_NVPTX_BCLIB_SUPPORTED})
       libomptarget_error_say("Cannot build CUDA LLVM bitcode offloading device RTL!")
     endif()
     libomptarget_say("Building CUDA LLVM bitcode offloading device RTL.")
 
     # Set flags for LLVM Bitcode compilation.
-    set(bc_flags ${LIBOMPTARGET_NVPTX_SELECTED_CUDA_COMPILER_FLAGS})
+    set(bc_flags ${LIBOMPTARGET_NVPTX_SELECTED_CUDA_COMPILER_FLAGS}
+                 -I${devicertl_base_directory})
     if(${LIBOMPTARGET_NVPTX_DEBUG})
       set(bc_flags ${bc_flags} -DOMPTARGET_NVPTX_DEBUG=-1)
     else()
       set(bc_flags ${bc_flags} -DOMPTARGET_NVPTX_DEBUG=0)
     endif()
 
     # CUDA 9 header files use the nv_weak attribute which clang is not yet prepared
     # to handle. Therefore, we use 'weak' instead. We are compiling only for the

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

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

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

 //===------------ option.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
 //
 //===----------------------------------------------------------------------===//
 //
 // GPU default options
 //
 //===----------------------------------------------------------------------===//
 #ifndef _OPTION_H_
 #define _OPTION_H_
 
+#include "interface.h"
+
 ////////////////////////////////////////////////////////////////////////////////
 // Kernel options
 ////////////////////////////////////////////////////////////////////////////////
 
 ////////////////////////////////////////////////////////////////////////////////
 // The following def must match the absolute limit hardwired in the host RTL
 // max number of threads per team
 #define MAX_THREADS_PER_TEAM 1024

 ////////////////////////////////////////////////////////////////////////////////
 // algo options
 ////////////////////////////////////////////////////////////////////////////////
 
 ////////////////////////////////////////////////////////////////////////////////
 // misc options (by def everythig here is device)
 ////////////////////////////////////////////////////////////////////////////////
 
-#define EXTERN extern "C" __device__
 #define INLINE __forceinline__ __device__
 #define NOINLINE __noinline__ __device__
 #ifndef TRUE
 #define TRUE 1
 #endif
 #ifndef FALSE
 #define FALSE 0
 #endif
 
 #endif

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

 }
 
 INLINE uint64_t __kmpc_impl_pack(uint32_t lo, uint32_t hi) {
   uint64_t val;
   asm volatile("mov.b64 %0, {%1,%2};" : "=l"(val) : "r"(lo), "r"(hi));
   return val;
 }
 
-typedef uint32_t __kmpc_impl_lanemask_t;
 static const __kmpc_impl_lanemask_t __kmpc_impl_all_lanes =
     UINT32_C(0xffffffff);
 
 INLINE __kmpc_impl_lanemask_t __kmpc_impl_lanemask_lt() {
   __kmpc_impl_lanemask_t res;
   asm("mov.u32 %0, %%lanemask_lt;" : "=r"(res));
   return res;
 }