diff --git a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp index 6a278fcdcf39e..a0cfa2167a161 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp @@ -1753,44 +1753,43 @@ void lowerGlobalsToVector(Module &M) { static void checkSLMInit(Module &M) { SmallPtrSet Callers; for (auto &F : M) { - if (isSlmInit(F)) { - for (User *U : F.users()) { - auto *FCall = dyn_cast(U); - if (FCall && FCall->getCalledFunction() == &F) { - Function *GenF = FCall->getFunction(); - SmallPtrSet Visited; - sycl::utils::traverseCallgraphUp( - GenF, - [&](Function *GraphNode) { - if (llvm::esimd::isESIMDKernel(*GraphNode)) { - if (Callers.contains(GraphNode)) { - StringRef KernelName = GraphNode->getName(); - std::string ErrorMsg = - std::string( - "slm_init is called more than once from kernel '") + - demangle(KernelName.str()) + "'."; - GraphNode->getContext().emitError(ErrorMsg); - } else { - Callers.insert(GraphNode); - } + if (!isSlmInit(F)) + continue; + for (User *U : F.users()) { + auto *FCall = dyn_cast(U); + if (FCall && FCall->getCalledFunction() == &F) { + Function *GenF = FCall->getFunction(); + SmallPtrSet Visited; + sycl::utils::traverseCallgraphUp( + GenF, + [&](Function *GraphNode) { + if (llvm::esimd::isESIMDKernel(*GraphNode)) { + if (Callers.contains(GraphNode)) { + StringRef KernelName = GraphNode->getName(); + std::string ErrorMsg = + std::string( + "slm_init is called more than once from kernel '") + + demangle(KernelName.str()) + "'."; + GraphNode->getContext().emitError(ErrorMsg); + } else { + Callers.insert(GraphNode); } - }, - Visited, false); - bool VisitedKernel = false; - for (const Function *Caller : Visited) { - if (llvm::esimd::isESIMDKernel(*Caller)) { - VisitedKernel = true; - break; - } - } - if (!VisitedKernel) { - F.getContext().emitError( - "slm_init must be called directly from ESIMD kernel."); + } + }, + Visited, false); + bool VisitedKernel = false; + for (const Function *Caller : Visited) { + if (llvm::esimd::isESIMDKernel(*Caller)) { + VisitedKernel = true; + break; } - } else { + } + if (!VisitedKernel) { F.getContext().emitError( - "slm_init can only be used as a direct call."); + "slm_init must be called directly from ESIMD kernel."); } + } else { + F.getContext().emitError("slm_init can only be used as a direct call."); } } }