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.
Details
Diff Detail
- Repository
- rG LLVM Github Monorepo
Unit Tests
Event Timeline
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 |
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.
openmp/libomptarget/plugins/cuda/src/rtl.cpp | ||
---|---|---|
1097 | Fixed it rG76c0c0ca86a1. |
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.
Nit, add a const/constexpr that semantically defines 2, e..g, const int SPMDGeneric =2