This is an archive of the discontinued LLVM Phabricator instance.

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

Authored by jhuber6 on Jul 21 2021, 10:01 AM.

Details

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.

Diff Detail

Event Timeline

jhuber6 created this revision.Jul 21 2021, 10:01 AM
jhuber6 requested review of this revision.Jul 21 2021, 10:01 AM
Herald added projects: Restricted Project, Restricted Project. · View Herald TranscriptJul 21 2021, 10:01 AM
ggeorgakoudis added inline comments.Jul 21 2021, 10:52 AM
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.

1971–1975

Nit, comment /* ExecutionMode == SPMD_GENERIC */

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

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

1096–1098

print explicitly "SPMD_GENERIC" when the mode is such

jhuber6 updated this revision to Diff 360530.Jul 21 2021, 11:04 AM

Addressing comments.

ggeorgakoudis accepted this revision.Jul 21 2021, 11:11 AM
This revision is now accepted and ready to land.Jul 21 2021, 11:11 AM
This revision was landed with ongoing or failed builds.Jul 21 2021, 5:57 PM
This revision was automatically updated to reflect the committed changes.
ye-luo added a subscriber: ye-luo.Jul 21 2021, 9:32 PM

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.

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 added inline comments.
openmp/libomptarget/plugins/cuda/src/rtl.cpp
1096–1098

There is a tailing space.

1097

This line exceeds character limits.

jhuber6 added inline comments.Jul 22 2021, 6:53 PM
openmp/libomptarget/plugins/cuda/src/rtl.cpp
1097

Fixed it rG76c0c0ca86a1.

tianshilei1992 added inline comments.Jul 22 2021, 7:08 PM
openmp/libomptarget/plugins/cuda/src/rtl.cpp
1096–1098

Thx

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.