[OpenMP] nvptx device RTL - final changes

[grokos]

Summary

This patch addresses the last few issues raised in the parent patch.

libomptarget/deviceRTLs/nvptx/CMakeLists.txt

if(LIBOMPTARGET_DEP_CUDA_FOUND)
   # that could be inlined in the user implementation.
   set(LIBOMPTARGET_NVPTX_ENABLE_BCLIB FALSE CACHE BOOL
     "Enable CUDA LLVM bitcode offloading device RTL.")
   if (${LIBOMPTARGET_NVPTX_ENABLE_BCLIB})
 
     # Find a clang compiler capable of compiling cuda files to LLVM bitcode and
     # an LLVM linker.
     # We use the one provided by the user, attempt to use the one used to build
-    # libomptarget, attempt to use clang in the PATH, or just fail.
+    # libomptarget or just fail.
 
     set(LIBOMPTARGET_NVPTX_CUDA_COMPILER "" CACHE STRING
       "Location of a CUDA compiler capable of emitting LLVM bitcode.")
     set(LIBOMPTARGET_NVPTX_BC_LINKER "" CACHE STRING
       "Location of a linker capable of linking LLVM bitcode objects.")
 
     if (NOT LIBOMPTARGET_NVPTX_CUDA_COMPILER STREQUAL "")
       set(LIBOMPTARGET_NVPTX_SELECTED_CUDA_COMPILER ${LIBOMPTARGET_NVPTX_CUDA_COMPILER})
     elseif(${CMAKE_C_COMPILER_ID} STREQUAL "Clang")
       set(LIBOMPTARGET_NVPTX_SELECTED_CUDA_COMPILER ${CMAKE_C_COMPILER})
     else()
-      libomptarget_say("Cannot find a CUDA compiler capable of emitting LLVM bitcode.")
-      libomptarget_say("Please configure with flag -DLIBOMPTARGET_NVPTX_CUDA_COMPILER")
+      libomptarget_error_say("Cannot find a CUDA compiler capable of emitting LLVM bitcode.")
+      libomptarget_error_say("Please configure with flag -DLIBOMPTARGET_NVPTX_CUDA_COMPILER")
     endif()
 
     # Get compiler directory to try to locate a suitable linker
     get_filename_component(COMPILER_DIR ${CMAKE_C_COMPILER} DIRECTORY)
     
     if (NOT LIBOMPTARGET_NVPTX_BC_LINKER STREQUAL "")
       set(LIBOMPTARGET_NVPTX_SELECTED_BC_LINKER ${LIBOMPTARGET_NVPTX_BC_LINKER})
     elseif(${CMAKE_C_COMPILER_ID} STREQUAL "Clang" AND EXISTS "${COMPILER_DIR}/llvm-link")
       # Use llvm-link from the directory containing clang
       set(LIBOMPTARGET_NVPTX_SELECTED_BC_LINKER ${COMPILER_DIR}/llvm-link)
     else()
-      libomptarget_say("Cannot find a linker capable of linking LLVM bitcode objects.")
-      libomptarget_say("Please configure with flag -DLIBOMPTARGET_NVPTX_BC_LINKER")
+      libomptarget_error_say("Cannot find a linker capable of linking LLVM bitcode objects.")
+      libomptarget_error_say("Please configure with flag -DLIBOMPTARGET_NVPTX_BC_LINKER")
     endif()
 
     if(LIBOMPTARGET_NVPTX_SELECTED_CUDA_COMPILER AND LIBOMPTARGET_NVPTX_SELECTED_BC_LINKER)
       libomptarget_say("Building CUDA LLVM bitcode offloading device RTL.")
 
       # Decide which ptx version to use. Same choices as Clang.
       if(CUDA_VERSION_MAJOR GREATER 9 OR CUDA_VERSION_MAJOR EQUAL 9)
         set(CUDA_PTX_VERSION ptx60)

libomptarget/deviceRTLs/nvptx/src/data_sharing.cu

__device__ static unsigned getNumThreads() {
   return blockDim.x;
 }
 // Thread ID in the CUDA block
 __device__ static unsigned getThreadId() {
   return threadIdx.x;
 }
 // Warp ID in the CUDA block
 __device__ static unsigned getWarpId() {
-  return threadIdx.x >> DS_Max_Worker_Warp_Size_Bits;
+  return threadIdx.x / WARPSIZE;
 }
 
 // The CUDA thread ID of the master thread.
 __device__ static unsigned getMasterThreadId() {
   unsigned Mask = WARPSIZE - 1;
   return (getNumThreads() - 1) & (~Mask);
 }
 
EXTERN void* __kmpc_get_data_sharing_environment_frame(int32_t SourceThreadID,
 
   // If the runtime has been elided, use __shared__ memory for master-worker
   // data sharing.  We're reusing the statically allocated data structure
   // that is used for standard data sharing.
   if (!IsOMPRuntimeInitialized) return (void *) &DataSharingState;
 
   // Get the frame used by the requested thread.
 
-  unsigned SourceWID = SourceThreadID >> DS_Max_Worker_Warp_Size_Bits;
+  unsigned SourceWID = SourceThreadID / WARPSIZE;
 
   DSPRINT(DSFLAG,"Source  warp: %d\n", SourceWID);
 
   void *P = DataSharingState.FramePtr[SourceWID];
   DSPRINT0(DSFLAG,"Exiting __kmpc_get_data_sharing_environment_frame\n");
   return P;
 }

libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.cu

 ////////////////////////////////////////////////////////////////////////////////
 
 INLINE unsigned smid() {
   unsigned id;
   asm("mov.u32 %0, %%smid;" : "=r"(id));
   return id;
 }
 
-INLINE unsigned n_sm() {
-  unsigned n_sm;
-  asm("mov.u32 %0, %%nsmid;" : "=r"(n_sm));
-  return n_sm;
-}
-
 EXTERN void __kmpc_kernel_init_params(void *Ptr) {
   PRINT(LD_IO, "call to __kmpc_kernel_init_params with version %f\n",
         OMPTARGET_NVPTX_VERSION);
 
   SetTeamsReductionScratchpadPtr(Ptr);
 }
 
 EXTERN void __kmpc_kernel_init(int ThreadLimit,
EXTERN void __kmpc_spmd_kernel_init(int ThreadLimit,
   workDescr.CounterGroup().Init(
       omptarget_nvptx_threadPrivateContext->Priv(threadId));
   PRINT(LD_PAR, "thread will execute parallel region with id %d in a team of "
                 "%d threads\n",
                 newTaskDescr->ThreadId(), newTaskDescr->NThreads());
 
   if (RequiresDataSharing && threadId % WARPSIZE == 0) {
     // Warp master innitializes data sharing environment.
-    unsigned WID = threadId >> DS_Max_Worker_Warp_Size_Bits;
+    unsigned WID = threadId / WARPSIZE;
     __kmpc_data_sharing_slot *RootS = currTeamDescr.RootS(WID);
     DataSharingState.SlotPtr[WID] = RootS;
     DataSharingState.StackPtr[WID] = (void*)&RootS->Data[0];
   }
 }
 
 EXTERN void __kmpc_spmd_kernel_deinit() {
   // We're not going to pop the task descr stack of each thread since

libomptarget/deviceRTLs/nvptx/src/option.h

 #define _OPTION_H_
 
 ////////////////////////////////////////////////////////////////////////////////
 // Kernel options
 ////////////////////////////////////////////////////////////////////////////////
 
 ////////////////////////////////////////////////////////////////////////////////
 // The following def must match the absolute limit hardwired in the host RTL
-#define THREAD_ABSOLUTE_LIMIT 1024 /* omptx limit (must match threadAbsoluteLimit) */
-
 // max number of threads per team
-#define MAX_THREADS_PER_TEAM THREAD_ABSOLUTE_LIMIT
+#define MAX_THREADS_PER_TEAM 1024
 
 #define WARPSIZE 32
 
 // The named barrier for active parallel threads of a team in an L1 parallel region
 // to synchronize with each other.
 #define L1_BARRIER (1)
 
 // Maximum number of omp state objects per SM allocated statically in global memory.

libomptarget/deviceRTLs/nvptx/src/reduction.cu

if (NumThreads > WARPSIZE) {
     return BlockThreadId == 0;
   } else if (isRuntimeUninitialized /* Never an L2 parallel region without the OMP runtime */) {
     return BlockThreadId == 0;
   }
 
   // Get the OMP thread Id. This is different from BlockThreadId in the case of
   // an L2 parallel region.
   return GetOmpThreadId(BlockThreadId, isSPMDExecutionMode, isRuntimeUninitialized) == 0;
-}
 #endif // __CUDA_ARCH__ >= 700
+}
 
 EXTERN
 int32_t __kmpc_nvptx_parallel_reduce_nowait(
     int32_t global_tid, int32_t num_vars, size_t reduce_size, void *reduce_data,
     kmp_ShuffleReductFctPtr shflFct, kmp_InterWarpCopyFctPtr cpyFct) {
   return nvptx_parallel_reduce_nowait(global_tid, num_vars, reduce_size,
                                       reduce_data, shflFct, cpyFct,
                                       /*isSPMDExecutionMode=*/isSPMDMode());
if (ActiveThreads > WARPSIZE) {
     // Gather all the reduced values from each warp
     // to the first warp.
     cpyFct(reduce_data, WarpsNeeded);
 
     uint32_t WarpId = ThreadId/WARPSIZE;
     if (WarpId == 0)
       gpu_irregular_warp_reduce(reduce_data, shflFct, WarpsNeeded, ThreadId);
   }
+#endif // __CUDA_ARCH__ >= 700
 
   return ThreadId == 0;
 }
-#endif // __CUDA_ARCH__ >= 700
 
 EXTERN
 int32_t __kmpc_nvptx_teams_reduce_nowait(int32_t global_tid, int32_t num_vars,
                                          size_t reduce_size, void *reduce_data,
                                          kmp_ShuffleReductFctPtr shflFct,
                                          kmp_InterWarpCopyFctPtr cpyFct,
                                          kmp_CopyToScratchpadFctPtr scratchFct,
                                          kmp_LoadReduceFctPtr ldFct) {