@@ -73,7 +73,7 @@ sycl::getKernelNamesUsingImplicitLocalMem(const Module &M) {
7373 llvm::for_each (M.functions (), [&](const Function &F) {
7474 if (F.getCallingConv () == CallingConv::SPIR_KERNEL) {
7575 int ArgPos = GetArgumentPos (F);
76- if (ArgPos >= 0 )
76+ if (ArgPos >= 0 || F. hasFnAttribute (WORK_GROUP_STATIC_ATTR) )
7777 SPIRKernelNames.emplace_back (F.getName (), ArgPos);
7878 }
7979 });
@@ -190,6 +190,12 @@ static void lowerLocalMemCall(Function *LocalMemAllocFunc,
190190 auto *CI = cast<CallInst>(U);
191191 TransformCall (CI);
192192 DelCalls.push_back (CI);
193+ // Now, take each kernel that calls the builtins that allocate local memory,
194+ // either directly or through a series of function calls that eventually end
195+ // up in a direct call to the builtin, and attach the
196+ // work-group-memory-static attribute to the kernel if not already attached.
197+ // This is needed because free function kernels do not have the attribute
198+ // added by the library as is the case with other types of kernels.
193199 if (!FuncsCache.insert (CI->getFunction ()).second )
194200 continue ; // We have already traversed call graph from this function
195201
@@ -199,7 +205,7 @@ static void lowerLocalMemCall(Function *LocalMemAllocFunc,
199205 auto *F = WorkList.back ();
200206 WorkList.pop_back ();
201207
202- // Mark kernel as using scrach memory if it isn't marked already
208+ // Mark kernel as using scratch memory if it isn't marked already
203209 if (F->getCallingConv () == CallingConv::SPIR_KERNEL &&
204210 !F->hasFnAttribute (WORK_GROUP_STATIC_ATTR))
205211 F->addFnAttr (WORK_GROUP_STATIC_ATTR);
0 commit comments