Skip to content

Commit 764c842

Browse files
JonChesterfieldJonChesterfield
committed
[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
1 parent a2240f5 commit 764c842

File tree

6 files changed

+58
-35
lines changed

6 files changed

+58
-35
lines changed

‎openmp/libomptarget/deviceRTLs/nvptx/CMakeLists.txt

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -55,6 +55,7 @@ if(LIBOMPTARGET_DEP_CUDA_FOUND)
5555
src/omptarget-nvptx.cu
5656
src/parallel.cu
5757
src/reduction.cu
58+
src/support.cu
5859
src/sync.cu
5960
src/task.cu
6061
)

‎openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -385,6 +385,5 @@ INLINE omptarget_nvptx_TaskDescr *getMyTopTaskDescriptor(int globalThreadId);
385385
////////////////////////////////////////////////////////////////////////////////
386386

387387
#include "omptarget-nvptxi.h"
388-
#include "supporti.h"
389388

390389
#endif

‎openmp/libomptarget/deviceRTLs/nvptx/src/supporti.h renamed to ‎openmp/libomptarget/deviceRTLs/nvptx/src/support.cu

Lines changed: 5 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
//===--------- supporti.h - NVPTX OpenMP support functions ------- CUDA -*-===//
1+
//===--------- support.cu - NVPTX OpenMP support functions ------- CUDA -*-===//
22
//
33
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
44
// See https://llvm.org/LICENSE.txt for license information.
@@ -10,12 +10,14 @@
1010
//
1111
//===----------------------------------------------------------------------===//
1212

13+
#include "support.h"
14+
#include "debug.h"
15+
#include "omptarget-nvptx.h"
16+
1317
////////////////////////////////////////////////////////////////////////////////
1418
// Execution Parameters
1519
////////////////////////////////////////////////////////////////////////////////
1620

17-
#include "target_impl.h"
18-
1921
INLINE void setExecutionParameters(ExecutionMode EMode, RuntimeMode RMode) {
2022
execution_param = EMode;
2123
execution_param |= RMode;

‎openmp/libomptarget/deviceRTLs/nvptx/src/support.h

Lines changed: 48 additions & 29 deletions
Original file line numberDiff line numberDiff line change
@@ -10,7 +10,12 @@
1010
//
1111
//===----------------------------------------------------------------------===//
1212

13+
#ifndef OMPTARGET_SUPPORT_H
14+
#define OMPTARGET_SUPPORT_H
15+
16+
#include "interface.h"
1317
#include "target_impl.h"
18+
1419
////////////////////////////////////////////////////////////////////////////////
1520
// Execution Parameters
1621
////////////////////////////////////////////////////////////////////////////////
@@ -26,58 +31,70 @@ enum RuntimeMode {
2631
RuntimeMask = 0x02u,
2732
};
2833

29-
INLINE void setExecutionParameters(ExecutionMode EMode, RuntimeMode RMode);
30-
INLINE bool isGenericMode();
31-
INLINE bool isSPMDMode();
32-
INLINE bool isRuntimeUninitialized();
33-
INLINE bool isRuntimeInitialized();
34+
DEVICE void setExecutionParameters(ExecutionMode EMode, RuntimeMode RMode);
35+
DEVICE bool isGenericMode();
36+
DEVICE bool isSPMDMode();
37+
DEVICE bool isRuntimeUninitialized();
38+
DEVICE bool isRuntimeInitialized();
39+
40+
////////////////////////////////////////////////////////////////////////////////
41+
// Execution Modes based on location parameter fields
42+
////////////////////////////////////////////////////////////////////////////////
43+
44+
DEVICE bool checkSPMDMode(kmp_Ident *loc);
45+
46+
DEVICE bool checkGenericMode(kmp_Ident *loc);
47+
48+
DEVICE bool checkRuntimeUninitialized(kmp_Ident *loc);
49+
50+
DEVICE bool checkRuntimeInitialized(kmp_Ident *loc);
3451

3552
////////////////////////////////////////////////////////////////////////////////
3653
// get info from machine
3754
////////////////////////////////////////////////////////////////////////////////
3855

3956
// get low level ids of resources
40-
INLINE int GetThreadIdInBlock();
41-
INLINE int GetBlockIdInKernel();
42-
INLINE int GetNumberOfBlocksInKernel();
43-
INLINE int GetNumberOfThreadsInBlock();
44-
INLINE unsigned GetWarpId();
45-
INLINE unsigned GetLaneId();
57+
DEVICE int GetThreadIdInBlock();
58+
DEVICE int GetBlockIdInKernel();
59+
DEVICE int GetNumberOfBlocksInKernel();
60+
DEVICE int GetNumberOfThreadsInBlock();
61+
DEVICE unsigned GetWarpId();
62+
DEVICE unsigned GetLaneId();
4663

4764
// get global ids to locate tread/team info (constant regardless of OMP)
48-
INLINE int GetLogicalThreadIdInBlock(bool isSPMDExecutionMode);
49-
INLINE int GetMasterThreadID();
50-
INLINE int GetNumberOfWorkersInTeam();
65+
DEVICE int GetLogicalThreadIdInBlock(bool isSPMDExecutionMode);
66+
DEVICE int GetMasterThreadID();
67+
DEVICE int GetNumberOfWorkersInTeam();
5168

5269
// get OpenMP thread and team ids
53-
INLINE int GetOmpThreadId(int threadId,
70+
DEVICE int GetOmpThreadId(int threadId,
5471
bool isSPMDExecutionMode); // omp_thread_num
55-
INLINE int GetOmpTeamId(); // omp_team_num
72+
DEVICE int GetOmpTeamId(); // omp_team_num
5673

5774
// get OpenMP number of threads and team
58-
INLINE int GetNumberOfOmpThreads(bool isSPMDExecutionMode); // omp_num_threads
59-
INLINE int GetNumberOfOmpTeams(); // omp_num_teams
75+
DEVICE int GetNumberOfOmpThreads(bool isSPMDExecutionMode); // omp_num_threads
76+
DEVICE int GetNumberOfOmpTeams(); // omp_num_teams
6077

6178
// get OpenMP number of procs
62-
INLINE int GetNumberOfProcsInTeam(bool isSPMDExecutionMode);
63-
INLINE int GetNumberOfProcsInDevice(bool isSPMDExecutionMode);
79+
DEVICE int GetNumberOfProcsInTeam(bool isSPMDExecutionMode);
80+
DEVICE int GetNumberOfProcsInDevice(bool isSPMDExecutionMode);
6481

6582
// masters
66-
INLINE int IsTeamMaster(int ompThreadId);
83+
DEVICE int IsTeamMaster(int ompThreadId);
6784

6885
// Parallel level
69-
INLINE void IncParallelLevel(bool ActiveParallel, __kmpc_impl_lanemask_t Mask);
70-
INLINE void DecParallelLevel(bool ActiveParallel, __kmpc_impl_lanemask_t Mask);
86+
DEVICE void IncParallelLevel(bool ActiveParallel, __kmpc_impl_lanemask_t Mask);
87+
DEVICE void DecParallelLevel(bool ActiveParallel, __kmpc_impl_lanemask_t Mask);
7188

7289
////////////////////////////////////////////////////////////////////////////////
7390
// Memory
7491
////////////////////////////////////////////////////////////////////////////////
7592

7693
// safe alloc and free
77-
INLINE void *SafeMalloc(size_t size, const char *msg); // check if success
78-
INLINE void *SafeFree(void *ptr, const char *msg);
94+
DEVICE void *SafeMalloc(size_t size, const char *msg); // check if success
95+
DEVICE void *SafeFree(void *ptr, const char *msg);
7996
// pad to a alignment (power of 2 only)
80-
INLINE unsigned long PadBytes(unsigned long size, unsigned long alignment);
97+
DEVICE unsigned long PadBytes(unsigned long size, unsigned long alignment);
8198
#define ADD_BYTES(_addr, _bytes) \
8299
((void *)((char *)((void *)(_addr)) + (_bytes)))
83100
#define SUB_BYTES(_addr, _bytes) \
@@ -86,6 +103,8 @@ INLINE unsigned long PadBytes(unsigned long size, unsigned long alignment);
86103
////////////////////////////////////////////////////////////////////////////////
87104
// Teams Reduction Scratchpad Helpers
88105
////////////////////////////////////////////////////////////////////////////////
89-
INLINE unsigned int *GetTeamsReductionTimestamp();
90-
INLINE char *GetTeamsReductionScratchpad();
91-
INLINE void SetTeamsReductionScratchpadPtr(void *ScratchpadPtr);
106+
DEVICE unsigned int *GetTeamsReductionTimestamp();
107+
DEVICE char *GetTeamsReductionScratchpad();
108+
DEVICE void SetTeamsReductionScratchpadPtr(void *ScratchpadPtr);
109+
110+
#endif

‎openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -15,8 +15,9 @@
1515
#include <cuda.h>
1616
#include "nvptx_interface.h"
1717

18-
#define INLINE __forceinline__ __device__
19-
#define NOINLINE __noinline__ __device__
18+
#define DEVICE __device__
19+
#define INLINE __forceinline__ DEVICE
20+
#define NOINLINE __noinline__ DEVICE
2021

2122
////////////////////////////////////////////////////////////////////////////////
2223
// Kernel options

‎openmp/libomptarget/deviceRTLs/nvptx/unity.cu

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -21,5 +21,6 @@
2121
#include "src/omptarget-nvptx.cu"
2222
#include "src/parallel.cu"
2323
#include "src/reduction.cu"
24+
#include "src/support.cu"
2425
#include "src/sync.cu"
2526
#include "src/task.cu"

0 commit comments

Comments
 (0)