[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
git-svn-id: https://llvm.org/svn/llvm-project/openmp/trunk@374919 91177308-0d34-0410-b5e6-96231b3b80d8
diff --git a/libomptarget/deviceRTLs/nvptx/src/interface.h b/libomptarget/deviceRTLs/interface.h
similarity index 99%
rename from libomptarget/deviceRTLs/nvptx/src/interface.h
rename to libomptarget/deviceRTLs/interface.h
index 4a84922..0f0f43e 100644
--- a/libomptarget/deviceRTLs/nvptx/src/interface.h
+++ b/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
diff --git a/libomptarget/deviceRTLs/nvptx/CMakeLists.txt b/libomptarget/deviceRTLs/nvptx/CMakeLists.txt
index c20339c..1cd13c5 100644
--- a/libomptarget/deviceRTLs/nvptx/CMakeLists.txt
+++ b/libomptarget/deviceRTLs/nvptx/CMakeLists.txt
@@ -35,6 +35,10 @@
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 @@
# 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 @@
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()
diff --git a/libomptarget/deviceRTLs/nvptx/src/nvptx_interface.h b/libomptarget/deviceRTLs/nvptx/src/nvptx_interface.h
new file mode 100644
index 0000000..7c9e471
--- /dev/null
+++ b/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
diff --git a/libomptarget/deviceRTLs/nvptx/src/option.h b/libomptarget/deviceRTLs/nvptx/src/option.h
index 0b04fdf..3c0beaf 100644
--- a/libomptarget/deviceRTLs/nvptx/src/option.h
+++ b/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
diff --git a/libomptarget/deviceRTLs/nvptx/src/target_impl.h b/libomptarget/deviceRTLs/nvptx/src/target_impl.h
index 37a125d..de2776e 100644
--- a/libomptarget/deviceRTLs/nvptx/src/target_impl.h
+++ b/libomptarget/deviceRTLs/nvptx/src/target_impl.h
@@ -26,7 +26,6 @@
return val;
}
-typedef uint32_t __kmpc_impl_lanemask_t;
static const __kmpc_impl_lanemask_t __kmpc_impl_all_lanes =
UINT32_C(0xffffffff);
