2828
2929#include " MemoryManager.h"
3030
31+ #include " llvm/Frontend/OpenMP/OMPConstants.h"
32+
3133// Utility for retrieving and printing CUDA error string.
3234#ifdef OMPTARGET_DEBUG
3335#define CUDA_ERR_STRING (err ) \
@@ -71,28 +73,17 @@ struct FuncOrGblEntryTy {
7173 std::vector<__tgt_offload_entry> Entries;
7274};
7375
74- enum ExecutionModeType {
75- SPMD, // constructors, destructors,
76- // combined constructs (`teams distribute parallel for [simd]`)
77- GENERIC, // everything else
78- SPMD_GENERIC, // Generic kernel with SPMD execution
79- NONE
80- };
81-
8276// / Use a single entity to encode a kernel and a set of flags.
8377struct KernelTy {
8478 CUfunction Func;
8579
8680 // execution mode of kernel
87- // 0 - SPMD mode (without master warp)
88- // 1 - Generic mode (with master warp)
89- // 2 - SPMD mode execution with Generic mode semantics.
90- int8_t ExecutionMode;
81+ llvm::omp::OMPTgtExecModeFlags ExecutionMode;
9182
9283 // / Maximal number of threads per block for this kernel.
9384 int MaxThreadsPerBlock = 0 ;
9485
95- KernelTy (CUfunction _Func, int8_t _ExecutionMode)
86+ KernelTy (CUfunction _Func, llvm::omp::OMPTgtExecModeFlags _ExecutionMode)
9687 : Func(_Func), ExecutionMode(_ExecutionMode) {}
9788};
9889
@@ -867,7 +858,7 @@ class DeviceRTLTy {
867858 DPxPTR (E - HostBegin), E->name , DPxPTR (Func));
868859
869860 // default value GENERIC (in case symbol is missing from cubin file)
870- int8_t ExecModeVal = ExecutionModeType::GENERIC ;
861+ llvm::omp::OMPTgtExecModeFlags ExecModeVal ;
871862 std::string ExecModeNameStr (E->name );
872863 ExecModeNameStr += " _exec_mode"
873864 const char *ExecModeName = ExecModeNameStr.c_str ();
@@ -876,9 +867,9 @@ class DeviceRTLTy {
876867 size_t CUSize;
877868 Err = cuModuleGetGlobal (&ExecModePtr, &CUSize, Module, ExecModeName);
878869 if (Err == CUDA_SUCCESS) {
879- if (CUSize != sizeof (int8_t )) {
870+ if (CUSize != sizeof (llvm::omp::OMPTgtExecModeFlags )) {
880871 DP (" Loading global exec_mode '%s' - size mismatch (%zd != %zd)\n "
881- ExecModeName, CUSize, sizeof (int8_t ));
872+ ExecModeName, CUSize, sizeof (llvm::omp::OMPTgtExecModeFlags ));
882873 return nullptr ;
883874 }
884875
@@ -890,12 +881,6 @@ class DeviceRTLTy {
890881 CUDA_ERR_STRING (Err);
891882 return nullptr ;
892883 }
893-
894- if (ExecModeVal < 0 || ExecModeVal > 2 ) {
895- DP (" Error wrong exec_mode value specified in cubin file: %d\n "
896- ExecModeVal);
897- return nullptr ;
898- }
899884 } else {
900885 DP (" Loading global exec_mode '%s' - symbol missing, using default "
901886 " value GENERIC (1)\n "
@@ -1098,12 +1083,19 @@ class DeviceRTLTy {
10981083
10991084 KernelTy *KernelInfo = reinterpret_cast <KernelTy *>(TgtEntryPtr);
11001085
1086+ const bool IsSPMDGenericMode =
1087+ KernelInfo->ExecutionMode == llvm::omp::OMP_TGT_EXEC_MODE_GENERIC_SPMD;
1088+ const bool IsSPMDMode =
1089+ KernelInfo->ExecutionMode == llvm::omp::OMP_TGT_EXEC_MODE_SPMD;
1090+ const bool IsGenericMode =
1091+ KernelInfo->ExecutionMode == llvm::omp::OMP_TGT_EXEC_MODE_GENERIC;
1092+
11011093 int CudaThreadsPerBlock;
11021094 if (ThreadLimit > 0 ) {
11031095 DP (" Setting CUDA threads per block to requested %d\n "
11041096 CudaThreadsPerBlock = ThreadLimit;
11051097 // Add master warp if necessary
1106- if (KernelInfo-> ExecutionMode == GENERIC ) {
1098+ if (IsGenericMode ) {
11071099 DP (" Adding master warp: +%d threads\n " WarpSize );
11081100 CudaThreadsPerBlock += DeviceData[DeviceId].WarpSize ;
11091101 }
@@ -1136,13 +1128,21 @@ class DeviceRTLTy {
11361128 unsigned int CudaBlocksPerGrid;
11371129 if (TeamNum <= 0 ) {
11381130 if (LoopTripCount > 0 && EnvNumTeams < 0 ) {
1139- if (KernelInfo->ExecutionMode == SPMD) {
1131+ if (IsSPMDGenericMode) {
1132+ // If we reach this point, then we are executing a kernel that was
1133+ // transformed from Generic-mode to SPMD-mode. This kernel has
1134+ // SPMD-mode execution, but needs its blocks to be scheduled
1135+ // differently because the current loop trip count only applies to the
1136+ // `teams distribute` region and will create var too few blocks using
1137+ // the regular SPMD-mode method.
1138+ CudaBlocksPerGrid = LoopTripCount;
1139+ } else if (IsSPMDMode) {
11401140 // We have a combined construct, i.e. `target teams distribute
11411141 // parallel for [simd]`. We launch so many teams so that each thread
11421142 // will execute one iteration of the loop. round up to the nearest
11431143 // integer
11441144 CudaBlocksPerGrid = ((LoopTripCount - 1 ) / CudaThreadsPerBlock) + 1 ;
1145- } else if (KernelInfo-> ExecutionMode == GENERIC ) {
1145+ } else if (IsGenericMode ) {
11461146 // If we reach this point, then we have a non-combined construct, i.e.
11471147 // `teams distribute` with a nested `parallel for` and each team is
11481148 // assigned one iteration of the `distribute` loop. E.g.:
@@ -1156,16 +1156,9 @@ class DeviceRTLTy {
11561156 // Threads within a team will execute the iterations of the `parallel`
11571157 // loop.
11581158 CudaBlocksPerGrid = LoopTripCount;
1159- } else if (KernelInfo->ExecutionMode == SPMD_GENERIC) {
1160- // If we reach this point, then we are executing a kernel that was
1161- // transformed from Generic-mode to SPMD-mode. This kernel has
1162- // SPMD-mode execution, but needs its blocks to be scheduled
1163- // differently because the current loop trip count only applies to the
1164- // `teams distribute` region and will create var too few blocks using
1165- // the regular SPMD-mode method.
1166- CudaBlocksPerGrid = LoopTripCount;
11671159 } else {
1168- REPORT (" Unknown execution mode: %d\n " ExecutionMode );
1160+ REPORT (" Unknown execution mode: %d\n "
1161+ static_cast <int8_t >(KernelInfo->ExecutionMode ));
11691162 return OFFLOAD_FAIL;
11701163 }
11711164 DP (" Using %d teams due to loop trip count %"
@@ -1185,16 +1178,12 @@ class DeviceRTLTy {
11851178 }
11861179
11871180 INFO (OMP_INFOTYPE_PLUGIN_KERNEL, DeviceId,
1188- " Launching kernel %s with %d blocks and %d threads in %s "
1189- " mode\n "
1181+ " Launching kernel %s with %d blocks and %d threads in %s mode\n "
11901182 (getOffloadEntry (DeviceId, TgtEntryPtr))
11911183 ? getOffloadEntry (DeviceId, TgtEntryPtr)->name
11921184 : " (null)"
11931185 CudaBlocksPerGrid, CudaThreadsPerBlock,
1194- (KernelInfo->ExecutionMode != SPMD
1195- ? (KernelInfo->ExecutionMode == GENERIC ? " Generic"
1196- : " SPMD-Generic"
1197- : " SPMD"
1186+ (!IsSPMDMode ? (IsGenericMode ? " Generic" " SPMD-Generic" " SPMD"
11981187
11991188 CUstream Stream = getStream (DeviceId, AsyncInfo);
12001189 Err = cuLaunchKernel (KernelInfo->Func , CudaBlocksPerGrid, /* gridDimY */ 1 ,
0 commit comments