Page MenuHomePhabricator

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

Authored by jhuber6 on Wed, Jul 21, 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.Wed, Jul 21, 10:01 AM
jhuber6 requested review of this revision.Wed, Jul 21, 10:01 AM
Herald added projects: Restricted Project, Restricted Project. · View Herald TranscriptWed, Jul 21, 10:01 AM
ggeorgakoudis added inline comments.Wed, Jul 21, 10:52 AM
llvm/lib/Transforms/IPO/OpenMPOpt.cpp
2892

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

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

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

1970–1974

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.Wed, Jul 21, 11:04 AM

Addressing comments.

ggeorgakoudis accepted this revision.Wed, Jul 21, 11:11 AM
This revision is now accepted and ready to land.Wed, Jul 21, 11:11 AM
This revision was landed with ongoing or failed builds.Wed, Jul 21, 5:57 PM
This revision was automatically updated to reflect the committed changes.
ye-luo added a subscriber: ye-luo.Wed, Jul 21, 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

There is a tailing space.

1097

This line exceeds character limits.

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

Fixed it rG76c0c0ca86a1.

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

Thx