diff --git a/llvm/lib/SYCLLowerIR/LowerWGScope.cpp b/llvm/lib/SYCLLowerIR/LowerWGScope.cpp index 01d276ddaf662..0c107b07e3460 100644 --- a/llvm/lib/SYCLLowerIR/LowerWGScope.cpp +++ b/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(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) diff --git a/sycl/test-e2e/HierPar/hier_par_indirect.cpp b/sycl/test-e2e/HierPar/hier_par_indirect.cpp index 073013a48b63b..ae9dee5a25381 100644 --- a/sycl/test-e2e/HierPar/hier_par_indirect.cpp +++ b/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;