diff --git a/openmp/libomptarget/deviceRTLs/nvptx/CMakeLists.txt b/openmp/libomptarget/deviceRTLs/nvptx/CMakeLists.txt --- a/openmp/libomptarget/deviceRTLs/nvptx/CMakeLists.txt +++ b/openmp/libomptarget/deviceRTLs/nvptx/CMakeLists.txt @@ -55,6 +55,7 @@ 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 --- a/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h +++ b/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h @@ -385,6 +385,5 @@ //////////////////////////////////////////////////////////////////////////////// #include "omptarget-nvptxi.h" -#include "supporti.h" #endif diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/support.h b/openmp/libomptarget/deviceRTLs/nvptx/src/support.h --- 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 @@ 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 @@ //////////////////////////////////////////////////////////////////////////////// // 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/supporti.h b/openmp/libomptarget/deviceRTLs/nvptx/src/support.cu rename from openmp/libomptarget/deviceRTLs/nvptx/src/supporti.h rename to openmp/libomptarget/deviceRTLs/nvptx/src/support.cu --- 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/target_impl.h b/openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h --- a/openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h +++ b/openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h @@ -15,8 +15,9 @@ #include #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 --- 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"