[SYCL] Fix bug when using no device split and reqd_work_group_size

Author: jzc

llvm/include/llvm/SYCLLowerIR/ComputeModuleRuntimeInfo.h

@@ -34,7 +34,8 @@ using EntryPointSet = SetVector<Function *>;
 
 PropSetRegTy computeModuleProperties(const Module &M,
                                      const EntryPointSet &EntryPoints,
-                                     const GlobalBinImageProps &GlobProps);
+                                     const GlobalBinImageProps &GlobProps,
+                                     module_split::IRSplitMode SplitMode);
 
 std::string computeModuleSymbolTable(const Module &M,
                                      const EntryPointSet &EntryPoints);

llvm/lib/SYCLLowerIR/ComputeModuleRuntimeInfo.cpp

@@ -152,7 +152,8 @@ std::optional<T> getKernelSingleEltMetadata(const Function &Func,
 
 PropSetRegTy computeModuleProperties(const Module &M,
                                      const EntryPointSet &EntryPoints,
-                                     const GlobalBinImageProps &GlobProps) {
+                                     const GlobalBinImageProps &GlobProps,
+                                     module_split::IRSplitMode SplitMode) {
 
   PropSetRegTy PropSet;
   {
@@ -161,8 +162,16 @@ PropSetRegTy computeModuleProperties(const Module &M,
     PropSet.add(PropSetRegTy::SYCL_DEVICELIB_REQ_MASK, RMEntry);
   }
   {
+    // Usually, we would only expect one ReqdWGSize, as the module passed to
+    // this function would be split according to that. However, when splitting
+    // is disabled, this cannot be guaranteed. In this case, we reset the value,
+    // which makes so that no value is reqd_work_group_size data is attached in
+    // in the device image.
+    SYCLDeviceRequirements DeviceReqs = computeDeviceRequirements(M, EntryPoints);
+    if (SplitMode == module_split::SPLIT_NONE)
+      DeviceReqs.ReqdWorkGroupSize.reset();
     PropSet.add(PropSetRegTy::SYCL_DEVICE_REQUIREMENTS,
-                computeDeviceRequirements(M, EntryPoints).asMap());
+                DeviceReqs.asMap());
   }
 
   // extract spec constant maps per each module

llvm/lib/SYCLLowerIR/SYCLDeviceRequirements.cpp

@@ -40,7 +40,6 @@ SYCLDeviceRequirements
 llvm::computeDeviceRequirements(const Module &M,
                                 const SetVector<Function *> &EntryPoints) {
   SYCLDeviceRequirements Reqs;
-  bool MultipleReqdWGSize = false;
   // Process all functions in the module
   for (const Function &F : M) {
     if (auto *MDN = F.getMetadata("sycl_used_aspects")) {
@@ -81,8 +80,6 @@ llvm::computeDeviceRequirements(const Module &M,
             ExtractUnsignedIntegerFromMDNodeOperand(MDN, I));
       if (!Reqs.ReqdWorkGroupSize.has_value())
         Reqs.ReqdWorkGroupSize = NewReqdWorkGroupSize;
-      if (Reqs.ReqdWorkGroupSize != NewReqdWorkGroupSize)
-        MultipleReqdWGSize = true;
     }
 
     if (auto *MDN = F.getMetadata("sycl_joint_matrix")) {
@@ -119,13 +116,6 @@ llvm::computeDeviceRequirements(const Module &M,
     }
   }
 
-  // Usually, we would only expect one ReqdWGSize, as the module passed to
-  // this function would be split according to that. However, when splitting
-  // is disabled, this cannot be guaranteed. In this case, we reset the value,
-  // which makes so that no value is reqd_work_group_size data is attached in
-  // in the device image.
-  if (MultipleReqdWGSize)
-    Reqs.ReqdWorkGroupSize.reset();
   return Reqs;
 }
 

llvm/tools/sycl-post-link/sycl-post-link.cpp

@@ -308,7 +308,7 @@ std::string saveModuleProperties(module_split::ModuleDesc &MD,
                                  const GlobalBinImageProps &GlobProps, int I,
                                  StringRef Suff, StringRef Target = "") {
   auto PropSet =
-      computeModuleProperties(MD.getModule(), MD.entries(), GlobProps);
+      computeModuleProperties(MD.getModule(), MD.entries(), GlobProps, SplitMode);
 
   std::string NewSuff = Suff.str();
   if (!Target.empty()) {

sycl/test-e2e/Regression/no-split-reqd-wg-size-2.cpp

@@ -0,0 +1,18 @@
+// This test checks that with -fsycl-device-code-split=off, kernels
+// with different reqd_work_group_size dimensions can be launched.
+
+// RUN: %{build} -fsycl -fsycl-device-code-split=off -o %t.out
+// RUN: %{run} %t.out
+
+// UNSUPPORTED: hip
+
+#include <sycl/detail/core.hpp>
+
+using namespace sycl;
+
+int main(int argc, char **argv) {
+  queue q;
+  q.single_task([]{});
+  q.parallel_for(range<2>(24, 1), [=](auto) [[sycl::reqd_work_group_size(24,1)]] {});
+  return 0;
+}