-//===------- 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.
//
//===----------------------------------------------------------------------===//
//
-// 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
#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
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.")
# 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})
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()
--- /dev/null
+//===--- 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
#ifndef _OPTION_H_
#define _OPTION_H_
+#include "interface.h"
+
////////////////////////////////////////////////////////////////////////////////
// Kernel options
////////////////////////////////////////////////////////////////////////////////
// misc options (by def everythig here is device)
////////////////////////////////////////////////////////////////////////////////
-#define EXTERN extern "C" __device__
#define INLINE __forceinline__ __device__
#define NOINLINE __noinline__ __device__
#ifndef TRUE
return val;
}
-typedef uint32_t __kmpc_impl_lanemask_t;
static const __kmpc_impl_lanemask_t __kmpc_impl_all_lanes =
UINT32_C(0xffffffff);