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

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.

Reviewers: ABataev, jdoerfert, grokos, ronlieb, RaviNarayanaswamy

Reviewed By: jdoerfert

Subscribers: mgorny, openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D68615

llvm-svn: 374919

openmp/libomptarget/deviceRTLs/interface.h

@@ -1,4 +1,4 @@
-//===------- 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.
@@ -6,8 +6,6 @@
 //
 //===----------------------------------------------------------------------===//
 //
-// 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
@@ -18,8 +16,11 @@
 #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

openmp/libomptarget/deviceRTLs/nvptx/CMakeLists.txt

@@ -35,6 +35,10 @@ if(CUDA_HOST_COMPILER MATCHES clang)
   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.")
 
@@ -83,7 +87,7 @@ if(LIBOMPTARGET_DEP_CUDA_FOUND)
   # 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})
 
@@ -117,7 +121,8 @@ if(LIBOMPTARGET_DEP_CUDA_FOUND)
     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()

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

@@ -0,0 +1,17 @@
+//===--- 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

@@ -12,6 +12,8 @@
 #ifndef _OPTION_H_
 #define _OPTION_H_
 
+#include "interface.h"
+
 ////////////////////////////////////////////////////////////////////////////////
 // Kernel options
 ////////////////////////////////////////////////////////////////////////////////
@@ -54,7 +56,6 @@
 // misc options (by def everythig here is device)
 ////////////////////////////////////////////////////////////////////////////////
 
-#define EXTERN extern "C" __device__
 #define INLINE __forceinline__ __device__
 #define NOINLINE __noinline__ __device__
 #ifndef TRUE

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

@@ -26,7 +26,6 @@ INLINE uint64_t __kmpc_impl_pack(uint32_t lo, uint32_t hi) {
   return val;
 }
 
-typedef uint32_t __kmpc_impl_lanemask_t;
 static const __kmpc_impl_lanemask_t __kmpc_impl_all_lanes =
     UINT32_C(0xffffffff);