[OpenMP] Optimized default kernel launch parameters in CUDA plugin

[grokos]

Summary

This patch modifies the default target kernel launch parameters (num_teams and thread_limit). The default thread_limit is set to 128 threads per team. In SPMD mode the kernel is launched with 128 threads, in non-SPMD mode we use 96 threads (+32 of the master warp).

The default number of teams has been optimized as follows. For the constructs below:

#target teams distribute
#teams distribute
#target teams distribute simd
#teams distribute simd

if the associated loop trip count is N, then the kernel is launched with N teams.

[Hahnfeld]

Does this change result in a lower runtime? Last time I tested clang-ykt on Pascal GPUs, 1024 threads were really the best thing to do...

libomptarget/plugins/cuda/src/rtl.cpp

594–598

Just move this code under if (thread_limit > 0)?

622–624

So each block executes one iteration? What is left for the threads in each block?

[arpith-jacob]

Hi Jonas,

The numbers are based on my testing of the Rodinia benchmark on k40m.

We don't have a working compiler on Pascal as yet (many of the omptests fail on Pascal) so I have not benchmarked on that GPU. Our compiler exposes a bug in the CUDA toolkit that is being fixed. It is possible that 1024 threads perform better on Pascal (these are of course heuristics) so we should extend the functionality here once Pascal support is added to the compiler/runtime.

libomptarget/plugins/cuda/src/rtl.cpp

622–624

Correct.

This case is for the 'teams distribute' construct. The assumption is that there is a nested parallel construct in which the threads within the block participate. Example:

#pragma omp target teams distribute
for(...) {

#pragma omp parallel for reduction(..)
for(..) {}

}

[grokos]

Wrote inline comments to make clear what the new default launch configuration is about.

[Hahnfeld]

LGTM

libomptarget/plugins/cuda/src/rtl.cpp

public:
   // OpenMP Environment properties
   int EnvNumTeams;
   int EnvTeamLimit;
 
   //static int EnvNumThreads;
   static const int HardTeamLimit = 1<<16; // 64k
   static const int HardThreadLimit = 1024;
   static const int DefaultNumTeams = 128;
-  static const int DefaultNumThreads = 1024;
+  static const int DefaultNumThreads = 128;
 
   // Record entry point associated with device
   void addOffloadEntry(int32_t device_id, __tgt_offload_entry entry) {
     assert(device_id < (int32_t)FuncGblEntries.size() &&
            "Unexpected device id!");
     FuncOrGblEntryTy &E = FuncGblEntries[device_id];
 
     E.Entries.push_back(entry);
int32_t __tgt_rtl_run_target_team_region(int32_t device_id, void *tgt_entry_ptr,
 
   int cudaThreadsPerBlock;
 
   if (thread_limit > 0) {
     cudaThreadsPerBlock = thread_limit;
     DP("Setting CUDA threads per block to requested %d\n", thread_limit);
   } else {
     cudaThreadsPerBlock = DeviceInfo.NumThreads[device_id];
+    if (KernelInfo->ExecutionMode == GENERIC) {
+      // Leave room for the master warp which will be added below.
+      cudaThreadsPerBlock -= DeviceInfo.WarpSize[device_id];
+    }
     DP("Setting CUDA threads per block to default %d\n",
         DeviceInfo.NumThreads[device_id]);
   }
 
   // Add master warp if necessary
   if (KernelInfo->ExecutionMode == GENERIC) {
     cudaThreadsPerBlock += DeviceInfo.WarpSize[device_id];
     DP("Adding master warp: +%d threads\n", DeviceInfo.WarpSize[device_id]);
   }
 
   if (cudaThreadsPerBlock > DeviceInfo.ThreadsPerBlock[device_id]) {
     cudaThreadsPerBlock = DeviceInfo.ThreadsPerBlock[device_id];
     DP("Threads per block capped at device limit %d\n",
         DeviceInfo.ThreadsPerBlock[device_id]);
   }
 
   int kernel_limit;
   err = cuFuncGetAttribute(&kernel_limit,
       CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK, KernelInfo->Func);
   if (err == CUDA_SUCCESS) {
     if (kernel_limit < cudaThreadsPerBlock) {
       cudaThreadsPerBlock = kernel_limit;
       DP("Threads per block capped at kernel limit %d\n", kernel_limit);
     }
   }
 
   int cudaBlocksPerGrid;
   if (team_num <= 0) {
     if (loop_tripcount > 0 && DeviceInfo.EnvNumTeams < 0) {
+      if (KernelInfo->ExecutionMode == SPMD) {
         // round up to the nearest integer
         cudaBlocksPerGrid = ((loop_tripcount - 1) / cudaThreadsPerBlock) + 1;
+      } else {
+        cudaBlocksPerGrid = loop_tripcount;
+      }
       DP("Using %d teams due to loop trip count %" PRIu64 " and number of "
           "threads per block %d\n", cudaBlocksPerGrid, loop_tripcount,
           cudaThreadsPerBlock);
     } else {
       cudaBlocksPerGrid = DeviceInfo.NumTeams[device_id];
       DP("Using default number of teams %d\n", DeviceInfo.NumTeams[device_id]);
     }
   } else if (team_num > DeviceInfo.BlocksPerGrid[device_id]) {