6 files changed, 58 insertions, 35 deletions
diff --git a/openmp/libomptarget/deviceRTLs/nvptx/CMakeLists.txt b/openmp/libomptarget/deviceRTLs/nvptx/CMakeLists.txt
index caf02d62c59..6e7f76f5c5b 100644
--- a/openmp/libomptarget/deviceRTLs/nvptx/CMakeLists.txt
+++ b/openmp/libomptarget/deviceRTLs/nvptx/CMakeLists.txt
@@ -55,6 +55,7 @@ if(LIBOMPTARGET_DEP_CUDA_FOUND)
       src/omptarget-nvptx.cu
       src/parallel.cu
       src/reduction.cu
+      src/support.cu
       src/sync.cu
       src/task.cu
   )
diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h b/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
index 986150402f1..6c98f0104ef 100644
--- a/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
+++ b/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
@@ -385,6 +385,5 @@ INLINE omptarget_nvptx_TaskDescr *getMyTopTaskDescriptor(int globalThreadId);
 ////////////////////////////////////////////////////////////////////////////////
 
 #include "omptarget-nvptxi.h"
-#include "supporti.h"
 
 #endif
diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/supporti.h b/openmp/libomptarget/deviceRTLs/nvptx/src/support.cu
index 6fa85789990..2b89d8d3466 100644
--- a/openmp/libomptarget/deviceRTLs/nvptx/src/supporti.h
+++ b/openmp/libomptarget/deviceRTLs/nvptx/src/support.cu
@@ -1,4 +1,4 @@
-//===--------- supporti.h - NVPTX OpenMP support functions ------- CUDA -*-===//
+//===--------- support.cu - NVPTX OpenMP 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.
@@ -10,12 +10,14 @@
 //
 //===----------------------------------------------------------------------===//
 
+#include "support.h"
+#include "debug.h"
+#include "omptarget-nvptx.h"
+
 ////////////////////////////////////////////////////////////////////////////////
 // Execution Parameters
 ////////////////////////////////////////////////////////////////////////////////
 
-#include "target_impl.h"
-
 INLINE void setExecutionParameters(ExecutionMode EMode, RuntimeMode RMode) {
   execution_param = EMode;
   execution_param |= RMode;
diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/support.h b/openmp/libomptarget/deviceRTLs/nvptx/src/support.h
index de685b89450..1adabaf67ea 100644
--- a/openmp/libomptarget/deviceRTLs/nvptx/src/support.h
+++ b/openmp/libomptarget/deviceRTLs/nvptx/src/support.h
@@ -10,7 +10,12 @@
 //
 //===----------------------------------------------------------------------===//
 
+#ifndef OMPTARGET_SUPPORT_H
+#define OMPTARGET_SUPPORT_H
+
+#include "interface.h"
 #include "target_impl.h"
+
 ////////////////////////////////////////////////////////////////////////////////
 // Execution Parameters
 ////////////////////////////////////////////////////////////////////////////////
@@ -26,58 +31,70 @@ enum RuntimeMode {
   RuntimeMask = 0x02u,
 };
 
-INLINE void setExecutionParameters(ExecutionMode EMode, RuntimeMode RMode);
-INLINE bool isGenericMode();
-INLINE bool isSPMDMode();
-INLINE bool isRuntimeUninitialized();
-INLINE bool isRuntimeInitialized();
+DEVICE void setExecutionParameters(ExecutionMode EMode, RuntimeMode RMode);
+DEVICE bool isGenericMode();
+DEVICE bool isSPMDMode();
+DEVICE bool isRuntimeUninitialized();
+DEVICE bool isRuntimeInitialized();
+
+////////////////////////////////////////////////////////////////////////////////
+// Execution Modes based on location parameter fields
+////////////////////////////////////////////////////////////////////////////////
+
+DEVICE bool checkSPMDMode(kmp_Ident *loc);
+
+DEVICE bool checkGenericMode(kmp_Ident *loc);
+
+DEVICE bool checkRuntimeUninitialized(kmp_Ident *loc);
+
+DEVICE bool checkRuntimeInitialized(kmp_Ident *loc);
 
 ////////////////////////////////////////////////////////////////////////////////
 // get info from machine
 ////////////////////////////////////////////////////////////////////////////////
 
 // get low level ids of resources
-INLINE int GetThreadIdInBlock();
-INLINE int GetBlockIdInKernel();
-INLINE int GetNumberOfBlocksInKernel();
-INLINE int GetNumberOfThreadsInBlock();
-INLINE unsigned GetWarpId();
-INLINE unsigned GetLaneId();
+DEVICE int GetThreadIdInBlock();
+DEVICE int GetBlockIdInKernel();
+DEVICE int GetNumberOfBlocksInKernel();
+DEVICE int GetNumberOfThreadsInBlock();
+DEVICE unsigned GetWarpId();
+DEVICE unsigned GetLaneId();
 
 // get global ids to locate tread/team info (constant regardless of OMP)
-INLINE int GetLogicalThreadIdInBlock(bool isSPMDExecutionMode);
-INLINE int GetMasterThreadID();
-INLINE int GetNumberOfWorkersInTeam();
+DEVICE int GetLogicalThreadIdInBlock(bool isSPMDExecutionMode);
+DEVICE int GetMasterThreadID();
+DEVICE int GetNumberOfWorkersInTeam();
 
 // get OpenMP thread and team ids
-INLINE int GetOmpThreadId(int threadId,
+DEVICE int GetOmpThreadId(int threadId,
                           bool isSPMDExecutionMode);    // omp_thread_num
-INLINE int GetOmpTeamId();                              // omp_team_num
+DEVICE int GetOmpTeamId();                              // omp_team_num
 
 // get OpenMP number of threads and team
-INLINE int GetNumberOfOmpThreads(bool isSPMDExecutionMode); // omp_num_threads
-INLINE int GetNumberOfOmpTeams();                           // omp_num_teams
+DEVICE int GetNumberOfOmpThreads(bool isSPMDExecutionMode); // omp_num_threads
+DEVICE int GetNumberOfOmpTeams();                           // omp_num_teams
 
 // get OpenMP number of procs
-INLINE int GetNumberOfProcsInTeam(bool isSPMDExecutionMode);
-INLINE int GetNumberOfProcsInDevice(bool isSPMDExecutionMode);
+DEVICE int GetNumberOfProcsInTeam(bool isSPMDExecutionMode);
+DEVICE int GetNumberOfProcsInDevice(bool isSPMDExecutionMode);
 
 // masters
-INLINE int IsTeamMaster(int ompThreadId);
+DEVICE int IsTeamMaster(int ompThreadId);
 
 // Parallel level
-INLINE void IncParallelLevel(bool ActiveParallel, __kmpc_impl_lanemask_t Mask);
-INLINE void DecParallelLevel(bool ActiveParallel, __kmpc_impl_lanemask_t Mask);
+DEVICE void IncParallelLevel(bool ActiveParallel, __kmpc_impl_lanemask_t Mask);
+DEVICE void DecParallelLevel(bool ActiveParallel, __kmpc_impl_lanemask_t Mask);
 
 ////////////////////////////////////////////////////////////////////////////////
 // Memory
 ////////////////////////////////////////////////////////////////////////////////
 
 // safe alloc and free
-INLINE void *SafeMalloc(size_t size, const char *msg); // check if success
-INLINE void *SafeFree(void *ptr, const char *msg);
+DEVICE void *SafeMalloc(size_t size, const char *msg); // check if success
+DEVICE void *SafeFree(void *ptr, const char *msg);
 // pad to a alignment (power of 2 only)
-INLINE unsigned long PadBytes(unsigned long size, unsigned long alignment);
+DEVICE unsigned long PadBytes(unsigned long size, unsigned long alignment);
 #define ADD_BYTES(_addr, _bytes)                                               \
   ((void *)((char *)((void *)(_addr)) + (_bytes)))
 #define SUB_BYTES(_addr, _bytes)                                               \
@@ -86,6 +103,8 @@ INLINE unsigned long PadBytes(unsigned long size, unsigned long alignment);
 ////////////////////////////////////////////////////////////////////////////////
 // Teams Reduction Scratchpad Helpers
 ////////////////////////////////////////////////////////////////////////////////
-INLINE unsigned int *GetTeamsReductionTimestamp();
-INLINE char *GetTeamsReductionScratchpad();
-INLINE void SetTeamsReductionScratchpadPtr(void *ScratchpadPtr);
+DEVICE unsigned int *GetTeamsReductionTimestamp();
+DEVICE char *GetTeamsReductionScratchpad();
+DEVICE void SetTeamsReductionScratchpadPtr(void *ScratchpadPtr);
+
+#endif
diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h b/openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h
index 95fe2ad3d3d..e0bacab3c64 100644
--- a/openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h
+++ b/openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h
@@ -15,8 +15,9 @@
 #include <cuda.h>
 #include "nvptx_interface.h"
 
-#define INLINE __forceinline__ __device__
-#define NOINLINE __noinline__ __device__
+#define DEVICE __device__
+#define INLINE __forceinline__ DEVICE
+#define NOINLINE __noinline__ DEVICE
 
 ////////////////////////////////////////////////////////////////////////////////
 // Kernel options
diff --git a/openmp/libomptarget/deviceRTLs/nvptx/unity.cu b/openmp/libomptarget/deviceRTLs/nvptx/unity.cu
index b6cfd0d8b9a..f5f92f3919e 100644
--- a/openmp/libomptarget/deviceRTLs/nvptx/unity.cu
+++ b/openmp/libomptarget/deviceRTLs/nvptx/unity.cu
@@ -21,5 +21,6 @@
 #include "src/omptarget-nvptx.cu"
 #include "src/parallel.cu"
 #include "src/reduction.cu"
+#include "src/support.cu"
 #include "src/sync.cu"
 #include "src/task.cu"