[SYCLLowerIR] Fix a hierarchical parallelism bug (#20550)

There is a bug in my earlier PR #20484. When encountering an indirect
call to a `parallel_for_work_item` function, the range of instructions
that will be subjected to the lowering pass must be closed, whereas the
current code simply skips over this instruction as its iterating the
instructions of the basic block which effectively keeps extending the
range. There is an E2E test that reflects this bug where it used to hang
before this fix but so far I'm not aware of an LLVM-IR snippet of code
that triggers it so for the moment, I'm just adding a SYCL-level test.

Author: lbushi25

llvm/lib/SYCLLowerIR/LowerWGScope.cpp

@@ -851,14 +851,22 @@ PreservedAnalyses SYCLLowerWGScopePass::run(Function &F,
         }
         continue;
       }
-      // In addition to an instruction not having side effects, we end the range
-      // if the instruction is a call that contains, possibly several layers
-      // down the stack, a call to a parallel_for_work_item. Such calls should
-      // not be subject to lowering since they must be executed by every work
-      // item.
+      // We also split the range if the instruction is a call that contains,
+      // possibly several layers down the stack, a call to a
+      // parallel_for_work_item. Such calls should not be subject to lowering
+      // since they must be executed by every work item.
       const CallInst *Call = dyn_cast<CallInst>(I);
-      if (!mayHaveSideEffects(I) ||
-          (Call && hasCallToAFuncWithPFWIMetadata(*Call->getCalledFunction())))
+      if (Call && hasCallToAFuncWithPFWIMetadata(*Call->getCalledFunction())) {
+        if (First) {
+          assert(Last && "range must have been closed 1");
+          Ranges.push_back(InstrRange{First, Last});
+          First = nullptr;
+          Last = nullptr;
+        }
+        continue;
+      }
+
+      if (!mayHaveSideEffects(I))
         continue;
       LLVM_DEBUG(llvm::dbgs() << "+++ Side effects: " << *I << "\n");
       if (!First)

sycl/test-e2e/HierPar/hier_par_indirect.cpp

@@ -48,6 +48,17 @@ int main(int argc, char **argv) {
            wGroup.parallel_for_work_item([&](sycl::h_item<1> index) {});
          }));
    }).wait();
+  // Also try an example of a work-group scope variable being used in work item
+  // scope
+  q.submit([&](sycl::handler &cgh) {
+     cgh.parallel_for_work_group(
+         sycl::range<1>{1}, sycl::range<1>{128}, ([=](sycl::group<1> wGroup) {
+           int data;
+           foo(wGroup);
+           wGroup.parallel_for_work_item(
+               [&](sycl::h_item<1> index) { data = 0; });
+         }));
+   }).wait();
 
   std::cout << "test passed" << std::endl;
   return 0;