[OpenMP] Add new execution mode for SPMD execution with Generic semantics

[jhuber6]

Summary

Qualified kernels can be transformed from generic-mode to SPMD mode using an
optimization in OpenMPOpt. This patch introduces a new execution mode to
indicate kernels that have been transformed from generic-mode to SPMD-mode.
These kernels have SPMD-mode execution, but need generic-mode semantics for
scheduling the blocks and threads. Without this far too few blocks will be
scheduled for a generic region as SPMD mode expects the trip count to be
divided by the number of threads.

[ggeorgakoudis]

llvm/lib/Transforms/IPO/OpenMPOpt.cpp

2860

Nit, add a const/constexpr that semantically defines 2, e..g, const int SPMDGeneric =2

openmp/libomptarget/plugins/amdgpu/src/rtl.cpp

134

Nit, rename SPMD_GENERIC or SPMDIZED to show affinity to SPMD. Other suggestions to naming welcome.

1973

Nit, comment /* ExecutionMode == SPMD_GENERIC */

openmp/libomptarget/plugins/cuda/src/rtl.cpp

1064

Nit, comment /* KernelInfo->ExecutionMode == SPMD_GENERIC */

1096

print explicitly "SPMD_GENERIC" when the mode is such

[jhuber6]

Addressing comments.

[ggeorgakoudis]

LGTM

[ye-luo]

Please fix the warning

/home/ci/ecp/working/jlselogin3/batch-03/intel-dga/automatedtests/yeluo/test_llvm/builds/users/yeluo/NFStdQBJ/0/intel-dga/automatedtests/yeluo/test_llvm/llvm-project/openmp/libomptarget/plugins/cuda/src/rtl.cpp:1064:20: warning: variable 'CudaBlocksPerGrid' is used uninitialized whenever 'if' condition is false [-Wsometimes-uninitialized]

} else if (KernelInfo->ExecutionMode == SPMD_GENERIC) {
           ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~

/home/ci/ecp/working/jlselogin3/batch-03/intel-dga/automatedtests/yeluo/test_llvm/builds/users/yeluo/NFStdQBJ/0/intel-dga/automatedtests/yeluo/test_llvm/llvm-project/openmp/libomptarget/plugins/cuda/src/rtl.cpp:1101:44: note: uninitialized use occurs here

Err = cuLaunchKernel(KernelInfo->Func, CudaBlocksPerGrid, /* gridDimY */ 1,
                                       ^~~~~~~~~~~~~~~~~

/home/ci/ecp/working/jlselogin3/batch-03/intel-dga/automatedtests/yeluo/test_llvm/builds/users/yeluo/NFStdQBJ/0/intel-dga/automatedtests/yeluo/test_llvm/llvm-project/openmp/libomptarget/plugins/cuda/src/rtl.cpp:1064:16: note: remove the 'if' if its condition is always true

} else if (KernelInfo->ExecutionMode == SPMD_GENERIC) {
       ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~

/home/ci/ecp/working/jlselogin3/batch-03/intel-dga/automatedtests/yeluo/test_llvm/builds/users/yeluo/NFStdQBJ/0/intel-dga/automatedtests/yeluo/test_llvm/llvm-project/openmp/libomptarget/plugins/cuda/src/rtl.cpp:1041:35: note: initialize the variable 'CudaBlocksPerGrid' to silence this warning

unsigned int CudaBlocksPerGrid;
                              ^
                               = 0

1 warning generated.

[jhuber6]

In D106460#2895470, @ye-luo wrote:

Please fix the warning

/home/ci/ecp/working/jlselogin3/batch-03/intel-dga/automatedtests/yeluo/test_llvm/builds/users/yeluo/NFStdQBJ/0/intel-dga/automatedtests/yeluo/test_llvm/llvm-project/openmp/libomptarget/plugins/cuda/src/rtl.cpp:1064:20: warning: variable 'CudaBlocksPerGrid' is used uninitialized whenever 'if' condition is false [-Wsometimes-uninitialized]

} else if (KernelInfo->ExecutionMode == SPMD_GENERIC) {
           ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~

/home/ci/ecp/working/jlselogin3/batch-03/intel-dga/automatedtests/yeluo/test_llvm/builds/users/yeluo/NFStdQBJ/0/intel-dga/automatedtests/yeluo/test_llvm/llvm-project/openmp/libomptarget/plugins/cuda/src/rtl.cpp:1101:44: note: uninitialized use occurs here

Err = cuLaunchKernel(KernelInfo->Func, CudaBlocksPerGrid, /* gridDimY */ 1,
                                       ^~~~~~~~~~~~~~~~~

/home/ci/ecp/working/jlselogin3/batch-03/intel-dga/automatedtests/yeluo/test_llvm/builds/users/yeluo/NFStdQBJ/0/intel-dga/automatedtests/yeluo/test_llvm/llvm-project/openmp/libomptarget/plugins/cuda/src/rtl.cpp:1064:16: note: remove the 'if' if its condition is always true

} else if (KernelInfo->ExecutionMode == SPMD_GENERIC) {
       ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~

/home/ci/ecp/working/jlselogin3/batch-03/intel-dga/automatedtests/yeluo/test_llvm/builds/users/yeluo/NFStdQBJ/0/intel-dga/automatedtests/yeluo/test_llvm/llvm-project/openmp/libomptarget/plugins/cuda/src/rtl.cpp:1041:35: note: initialize the variable 'CudaBlocksPerGrid' to silence this warning

unsigned int CudaBlocksPerGrid;
                              ^
                               = 0

1 warning generated.

rGa158d3663fc5

[tianshilei1992]

openmp/libomptarget/plugins/cuda/src/rtl.cpp

1096–1099

There is a tailing space.

1097

This line exceeds character limits.

[jhuber6]

openmp/libomptarget/plugins/cuda/src/rtl.cpp

1097

Fixed it rG76c0c0ca86a1.

[tianshilei1992]

openmp/libomptarget/plugins/cuda/src/rtl.cpp

1096

Thx

[JonChesterfield]

Hi! Thanks for updating the amdgpu plugin along with this work, it's taken the edge off working out what the semantics are meant to be. Please tag me as reviewer for future changes to it.

llvm/lib/Transforms/IPO/OpenMPOpt.cpp

bool changeToSPMDMode(Attributor &A) {
     // kernel is executed in.
     Function *Kernel = getAnchorScope();
     GlobalVariable *ExecMode = Kernel->getParent()->getGlobalVariable(
         (Kernel->getName() + "_exec_mode").str());
     assert(ExecMode && "Kernel without exec mode?");
     assert(ExecMode->getInitializer() &&
            ExecMode->getInitializer()->isOneValue() &&
            "Initially non-SPMD kernel has SPMD exec mode!");
+    // Set the global exec mode flag to indicate Generic-SPMD mode.
+    if (!ExecMode->getInitializer()->isZeroValue())
       ExecMode->setInitializer(
-        ConstantInt::get(ExecMode->getInitializer()->getType(), 0));
+          ConstantInt::get(ExecMode->getInitializer()->getType(), 2));
 
     // Next rewrite the init and deinit calls to indicate we use SPMD-mode now.
     const int InitIsSPMDArgNo = 1;
     const int DeinitIsSPMDArgNo = 1;
     const int InitUseStateMachineArgNo = 2;
 
     auto &Ctx = getAnchorValue().getContext();
     A.changeUseAfterManifest(KernelInitCB->getArgOperandUse(InitIsSPMDArgNo),

llvm/test/Transforms/OpenMP/is_spmd_exec_mode_fold.ll

 ; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: --function-signature --check-globals
 ; RUN: opt -S -passes=openmp-opt < %s | FileCheck %s
 target triple = "nvptx64"
 
 %struct.ident_t = type { i32, i32, i32, i32, i8* }
 
 @is_spmd_exec_mode = weak constant i8 0
 @will_be_spmd_exec_mode = weak constant i8 1
 @non_spmd_exec_mode = weak constant i8 1
 @will_not_be_spmd_exec_mode = weak constant i8 1
 @G = external global i8
 @llvm.compiler.used = appending global [4 x i8*] [i8* @is_spmd_exec_mode, i8* @will_be_spmd_exec_mode, i8* @non_spmd_exec_mode, i8* @will_not_be_spmd_exec_mode ], section "llvm.metadata"
 
 ;.
 ; CHECK: @[[IS_SPMD_EXEC_MODE:[a-zA-Z0-9_$"\\.-]+]] = weak constant i8 0
-; CHECK: @[[WILL_BE_SPMD_EXEC_MODE:[a-zA-Z0-9_$"\\.-]+]] = weak constant i8 0
+; CHECK: @[[WILL_BE_SPMD_EXEC_MODE:[a-zA-Z0-9_$"\\.-]+]] = weak constant i8 2
 ; CHECK: @[[NON_SPMD_EXEC_MODE:[a-zA-Z0-9_$"\\.-]+]] = weak constant i8 1
 ; CHECK: @[[WILL_NOT_BE_SPMD_EXEC_MODE:[a-zA-Z0-9_$"\\.-]+]] = weak constant i8 1
 ; CHECK: @[[G:[a-zA-Z0-9_$"\\.-]+]] = external global i8
 ; CHECK: @[[LLVM_COMPILER_USED:[a-zA-Z0-9_$"\\.-]+]] = appending global [4 x i8*] [i8* @is_spmd_exec_mode, i8* @will_be_spmd_exec_mode, i8* @non_spmd_exec_mode, i8* @will_not_be_spmd_exec_mode], section "llvm.metadata"
 ;.
 define weak void @is_spmd() {
 ; CHECK-LABEL: define {{[^@]+}}@is_spmd() {
 ; CHECK-NEXT:    [[I:%.*]] = call i32 @__kmpc_target_init(%struct.ident_t* null, i1 true, i1 false, i1 false)

llvm/test/Transforms/OpenMP/spmdization.ll

 @__omp_offloading_2c_38c77_sequential_loop_l4_exec_mode = weak constant i8 1
 @llvm.compiler.used = appending global [1 x i8*] [i8* @__omp_offloading_2c_38c77_sequential_loop_l4_exec_mode], section "llvm.metadata"
 
 ; The second argument of __kmpc_target_init and deinit is is set to true to indicate that we can run in SPMD mode.
 ; We also adjusted the global __omp_offloading_2c_38c77_sequential_loop_l4_exec_mode to have a zero initializer (which indicates SPMD mode to the runtime).
 ;.
 ; CHECK: @[[GLOB0:[0-9]+]] = private unnamed_addr constant [23 x i8] c"
 ; CHECK: @[[GLOB1:[0-9]+]] = private unnamed_addr constant [[STRUCT_IDENT_T:%.*]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* @[[GLOB0]], i32 0, i32 0) }, align 8
-; CHECK: @[[__OMP_OFFLOADING_2C_38C77_SEQUENTIAL_LOOP_L4_EXEC_MODE:[a-zA-Z0-9_$"\\.-]+]] = weak constant i8 0
+; CHECK: @[[__OMP_OFFLOADING_2C_38C77_SEQUENTIAL_LOOP_L4_EXEC_MODE:[a-zA-Z0-9_$"\\.-]+]] = weak constant i8 2
 ; CHECK: @[[LLVM_COMPILER_USED:[a-zA-Z0-9_$"\\.-]+]] = appending global [1 x i8*] [i8* @__omp_offloading_2c_38c77_sequential_loop_l4_exec_mode], section "llvm.metadata"
 ;.
 define weak void @__omp_offloading_2c_38c77_sequential_loop_l4() #0 {
 ; CHECK-LABEL: define {{[^@]+}}@__omp_offloading_2c_38c77_sequential_loop_l4
 ; CHECK-SAME: () #[[ATTR0:[0-9]+]] {
 ; CHECK-NEXT:  entry:
 ; CHECK-NEXT:    [[DOTZERO_ADDR:%.*]] = alloca i32, align 4
 ; CHECK-NEXT:    [[DOTTHREADID_TEMP_:%.*]] = alloca i32, align 4

openmp/libomptarget/plugins/amdgpu/src/rtl.cpp

 #include <assert.h>
 #include <cstdio>
 #include <cstdlib>
 #include <cstring>
 #include <elf.h>
 #include <fstream>
 #include <functional>
 #include <iostream>
 #include <libelf.h>
 #include <list>
 #include <memory>
 #include <mutex>
 #include <shared_mutex>
 #include <thread>
 #include <unordered_map>
 #include <vector>
 

 
 /// Keep entries table per device
 struct FuncOrGblEntryTy {
   __tgt_target_table Table;
   std::vector<__tgt_offload_entry> Entries;
 };
 
 enum ExecutionModeType {
-  SPMD,    // constructors, destructors,
+  SPMD,         // constructors, destructors,
                 // combined constructs (`teams distribute parallel for [simd]`)
-  GENERIC, // everything else
+  GENERIC,      // everything else
+  GENERIC_SPMD, // Generic kernel with SPMD execution
   NONE
 };
 
 struct KernelArgPool {
 private:
   static pthread_mutex_t mutex;
 
 public:

 std::unordered_map<std::string /*kernel*/, std::unique_ptr<KernelArgPool>>
     KernelArgPoolMap;
 
 /// Use a single entity to encode a kernel and a set of flags
 struct KernelTy {
   // execution mode of kernel
   // 0 - SPMD mode (without master warp)
   // 1 - Generic mode (with master warp)
+  // 2 - SPMD mode execution with Generic mode semantics.
   int8_t ExecutionMode;
   int16_t ConstWGSize;
   int32_t device_id;
   void *CallStackAddr = nullptr;
   const char *Name;
 
   KernelTy(int8_t _ExecutionMode, int16_t _ConstWGSize, int32_t _device_id,
            void *_CallStackAddr, const char *_Name,
if (err == HSA_STATUS_SUCCESS) {
           return NULL;
         }
 
         memcpy(&ExecModeVal, ExecModePtr, (size_t)varsize);
 
         DP("After loading global for %s ExecMode = %d\n", ExecModeName,
            ExecModeVal);
 
-        if (ExecModeVal < 0 || ExecModeVal > 1) {
+        if (ExecModeVal < 0 || ExecModeVal > 2) {
           DP("Error wrong exec_mode value specified in HSA code object file: "
              "%d\n",
              ExecModeVal);
           return NULL;
         }
       } else {
         DP("Loading global exec_mode '%s' - symbol missing, using default "
            "value "
if (Env.NumTeams > 0) {
     num_groups = (Env.TeamLimit < num_groups) ? Env.TeamLimit : num_groups;
     DP("Modifying teams based on Env.TeamLimit%d\n", Env.TeamLimit);
   } else {
     if (num_teams <= 0) {
       if (loop_tripcount > 0) {
         if (ExecutionMode == SPMD) {
           // round up to the nearest integer
           num_groups = ((loop_tripcount - 1) / threadsPerGroup) + 1;
+        } else if (ExecutionMode == GENERIC) {
+          num_groups = loop_tripcount;
         } else {
+          // This is a generic kernel that was transformed to use SPMD-mode
+          // execution but uses Generic-mode semantics for scheduling.
           num_groups = loop_tripcount;
         }
         DP("Using %d teams due to loop trip count %" PRIu64 " and number of "
            "threads per block %d\n",
            num_groups, loop_tripcount, threadsPerGroup);
       }
     } else {
       num_groups = num_teams;

openmp/libomptarget/plugins/cuda/src/rtl.cpp

 //===----RTLs/cuda/src/rtl.cpp - Target RTLs Implementation ------- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // RTL for CUDA machine
 //
 //===----------------------------------------------------------------------===//
 
 #include <cassert>
 #include <cstddef>
 #include <cuda.h>
 #include <list>
 #include <memory>
 #include <mutex>
 #include <string>
 #include <unordered_map>
 #include <vector>
 
 #include "Debug.h"

 
 /// Keep entries table per device.
 struct FuncOrGblEntryTy {
   __tgt_target_table Table;
   std::vector<__tgt_offload_entry> Entries;
 };
 
 enum ExecutionModeType {
-  SPMD, // constructors, destructors,
+  SPMD,         // constructors, destructors,
                 // combined constructs (`teams distribute parallel for [simd]`)
-  GENERIC, // everything else
+  GENERIC,      // everything else
+  GENERIC_SPMD, // Generic kernel with SPMD execution
   NONE
 };
 
 /// Use a single entity to encode a kernel and a set of flags.
 struct KernelTy {
   CUfunction Func;
 
   // execution mode of kernel
   // 0 - SPMD mode (without master warp)
   // 1 - Generic mode (with master warp)
+  // 2 - SPMD mode execution with Generic mode semantics.
   int8_t ExecutionMode;
 
   /// Maximal number of threads per block for this kernel.
   int MaxThreadsPerBlock = 0;
 
   KernelTy(CUfunction _Func, int8_t _ExecutionMode)
       : Func(_Func), ExecutionMode(_ExecutionMode) {}
 };
for (const __tgt_offload_entry *E = HostBegin; E != HostEnd; ++E) {
         if (Err != CUDA_SUCCESS) {
           REPORT("Error when copying data from device to host. Pointers: "
                  "host = " DPxMOD ", device = " DPxMOD ", size = %zd\n",
                  DPxPTR(&ExecModeVal), DPxPTR(ExecModePtr), CUSize);
           CUDA_ERR_STRING(Err);
           return nullptr;
         }
 
-        if (ExecModeVal < 0 || ExecModeVal > 1) {
+        if (ExecModeVal < 0 || ExecModeVal > 2) {
           DP("Error wrong exec_mode value specified in cubin file: %d\n",
              ExecModeVal);
           return nullptr;
         }
       } else {
         REPORT("Loading global exec_mode '%s' - symbol missing, using default "
                "value GENERIC (1)\n",
                ExecModeName);
int runTargetTeamRegion(const int DeviceId, void *TgtEntryPtr, void **TgtArgs,
     if (TeamNum <= 0) {
       if (LoopTripCount > 0 && EnvNumTeams < 0) {
         if (KernelInfo->ExecutionMode == SPMD) {
           // We have a combined construct, i.e. `target teams distribute
           // parallel for [simd]`. We launch so many teams so that each thread
           // will execute one iteration of the loop. round up to the nearest
           // integer
           CudaBlocksPerGrid = ((LoopTripCount - 1) / CudaThreadsPerBlock) + 1;
-        } else {
+        } else if (KernelInfo->ExecutionMode == GENERIC) {
           // If we reach this point, then we have a non-combined construct, i.e.
           // `teams distribute` with a nested `parallel for` and each team is
           // assigned one iteration of the `distribute` loop. E.g.:
           //
           // #pragma omp target teams distribute
           // for(...loop_tripcount...) {
           //   #pragma omp parallel for
           //   for(...) {}
           // }
           //
           // Threads within a team will execute the iterations of the `parallel`
           // loop.
           CudaBlocksPerGrid = LoopTripCount;
+        } else {
+          // If we reach this point, then we are executing a kernel that was
+          // transformed from Generic-mode to SPMD-mode. This kernel has
+          // SPMD-mode execution, but needs its blocks to be scheduled
+          // differently because the current loop trip count only applies to the
+          // `teams distribute` region and will create var too few blocks using
+          // the regular SPMD-mode method.
+          CudaBlocksPerGrid = LoopTripCount;
         }
         DP("Using %d teams due to loop trip count %" PRIu32
            " and number of threads per block %d\n",
            CudaBlocksPerGrid, LoopTripCount, CudaThreadsPerBlock);
       } else {
         DP("Using default number of teams %d\n", DeviceData[DeviceId].NumTeams);
         CudaBlocksPerGrid = DeviceData[DeviceId].NumTeams;
       }
     } else if (TeamNum > DeviceData[DeviceId].BlocksPerGrid) {
       DP("Capping number of teams to team limit %d\n",
          DeviceData[DeviceId].BlocksPerGrid);
       CudaBlocksPerGrid = DeviceData[DeviceId].BlocksPerGrid;
     } else {
       DP("Using requested number of teams %d\n", TeamNum);
       CudaBlocksPerGrid = TeamNum;
     }
 
     INFO(OMP_INFOTYPE_PLUGIN_KERNEL, DeviceId,
          "Launching kernel %s with %d blocks and %d threads in %s "
          "mode\n",
          (getOffloadEntry(DeviceId, TgtEntryPtr))
              ? getOffloadEntry(DeviceId, TgtEntryPtr)->name
              : "(null)",
          CudaBlocksPerGrid, CudaThreadsPerBlock,
-         (KernelInfo->ExecutionMode == SPMD) ? "SPMD" : "Generic");
+         (KernelInfo->ExecutionMode == SPMD ||
+          KernelInfo->ExecutionMode == GENERIC_SPMD)
+             ? "SPMD"
+             : "Generic");
 
     CUstream Stream = getStream(DeviceId, AsyncInfo);
     Err = cuLaunchKernel(KernelInfo->Func, CudaBlocksPerGrid, /* gridDimY */ 1,
                          /* gridDimZ */ 1, CudaThreadsPerBlock,
                          /* blockDimY */ 1, /* blockDimZ */ 1,
                          /* sharedMemBytes */ 0, Stream, &Args[0], nullptr);
     if (!checkResult(Err, "Error returned from cuLaunchKernel\n"))
       return OFFLOAD_FAIL;