[Openmp-commits] [openmp] 764c842 - [nfc][libomptarget] Reorganise support header
via Openmp-commits
openmp-commits at lists.llvm.org
Thu Oct 31 10:15:16 PDT 2019
Author: JonChesterfield
Date: 2019-10-31T17:15:02Z
New Revision: 764c8420e4b8fc11a9fa94d00f4ee617aa754cb2
URL: https://github.com/llvm/llvm-project/commit/764c8420e4b8fc11a9fa94d00f4ee617aa754cb2
DIFF: https://github.com/llvm/llvm-project/commit/764c8420e4b8fc11a9fa94d00f4ee617aa754cb2.diff
LOG: [nfc][libomptarget] Reorganise support header
Summary:
[nfc][libomptarget] Reorganise support header
All functions defined in support implementation are now declared in support.h
Reordered functions in support implementation to match the sequence in support.h
Added include guards to support.h
Added #include interface to support.h to provide kmp_Ident declaration
Move supporti.h to support.cu and s/INLINE/EXTERN/g
Add remaining includes to support.cu
A minor side effect is to change the name mangling of the support functions to
extern "C". If this matters another macro along the lines of INLINE/EXTERN
can be added - perhaps DEVICE as that's the obvious implementation.
Reviewers: jdoerfert, ABataev, grokos
Reviewed By: jdoerfert
Subscribers: mgorny, jfb, openmp-commits
Tags: #openmp
Differential Revision: https://reviews.llvm.org/D69652
Added:
openmp/libomptarget/deviceRTLs/nvptx/src/support.cu
Modified:
openmp/libomptarget/deviceRTLs/nvptx/CMakeLists.txt
openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
openmp/libomptarget/deviceRTLs/nvptx/src/support.h
openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h
openmp/libomptarget/deviceRTLs/nvptx/unity.cu
Removed:
openmp/libomptarget/deviceRTLs/nvptx/src/supporti.h
################################################################################
diff --git a/openmp/libomptarget/deviceRTLs/nvptx/CMakeLists.txt b/openmp/libomptarget/deviceRTLs/nvptx/CMakeLists.txt
index caf02d62c59d..6e7f76f5c5b8 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 986150402f16..6c98f0104efc 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
similarity index 98%
rename from openmp/libomptarget/deviceRTLs/nvptx/src/supporti.h
rename to openmp/libomptarget/deviceRTLs/nvptx/src/support.cu
index 6fa857899905..2b89d8d34660 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 de685b894509..1adabaf67ea7 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 95fe2ad3d3d5..e0bacab3c643 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 b6cfd0d8b9a8..f5f92f3919e7 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"
More information about the Openmp-commits
mailing list