[SYCL][NFC] Helper for implicit kernarg targets in LowerWGLocalMemory (#16766)

To make the code more explicit and allow for future expansion.

Author: Maetveis

llvm/lib/SYCLLowerIR/LowerWGLocalMemory.cpp

@@ -49,12 +49,16 @@ class SYCLLowerWGLocalMemoryLegacy : public ModulePass {
 };
 } // namespace
 
+static bool usesKernelArgForDynWGLocalMem(const Triple &TT) {
+  return TT.isSPIROrSPIRV();
+}
+
 std::vector<std::pair<StringRef, int>>
 sycl::getKernelNamesUsingImplicitLocalMem(const Module &M) {
   std::vector<std::pair<StringRef, int>> SPIRKernelNames;
   Triple TT(M.getTargetTriple());
 
-  if (TT.isSPIROrSPIRV()) {
+  if (usesKernelArgForDynWGLocalMem(TT)) {
     auto GetArgumentPos = [&](const Function &F) -> int {
       for (const Argument &Arg : F.args())
         if (F.getAttributes().hasParamAttr(Arg.getArgNo(),
@@ -129,7 +133,7 @@ lowerDynamicLocalMemCallDirect(CallInst *CI, Triple TT,
 
   Value *GVPtr = [&]() -> Value * {
     IRBuilder<> Builder(CI);
-    if (TT.isSPIROrSPIRV())
+    if (usesKernelArgForDynWGLocalMem(TT))
       return Builder.CreateLoad(CI->getType(), LocalMemPlaceholder);
 
     return Builder.CreatePointerCast(LocalMemPlaceholder, CI->getType());
@@ -188,32 +192,33 @@ static bool dynamicWGLocalMemory(Module &M) {
   if (!LocalMemArrayGV) {
     assert(DLMFunc->isDeclaration() && "should have declaration only");
     Type *LocalMemArrayTy =
-        TT.isSPIROrSPIRV()
+        usesKernelArgForDynWGLocalMem(TT)
             ? static_cast<Type *>(PointerType::get(M.getContext(), LocalAS))
             : static_cast<Type *>(
                   ArrayType::get(Type::getInt8Ty(M.getContext()), 0));
     LocalMemArrayGV = new GlobalVariable(
         M,               // module
         LocalMemArrayTy, // type
         false,           // isConstant
-        TT.isSPIROrSPIRV() ? GlobalValue::LinkOnceODRLinkage
-                           : GlobalValue::ExternalLinkage, // Linkage
-        TT.isSPIROrSPIRV() ? UndefValue::get(LocalMemArrayTy)
-                           : nullptr,   // Initializer
-        DYNAMIC_LOCALMEM_GV,            // Name prefix
-        nullptr,                        // InsertBefore
-        GlobalVariable::NotThreadLocal, // ThreadLocalMode
-        LocalAS                         // AddressSpace
+        usesKernelArgForDynWGLocalMem(TT)
+            ? GlobalValue::LinkOnceODRLinkage
+            : GlobalValue::ExternalLinkage, // Linkage
+        usesKernelArgForDynWGLocalMem(TT) ? UndefValue::get(LocalMemArrayTy)
+                                          : nullptr, // Initializer
+        DYNAMIC_LOCALMEM_GV,                         // Name prefix
+        nullptr,                                     // InsertBefore
+        GlobalVariable::NotThreadLocal,              // ThreadLocalMode
+        LocalAS                                      // AddressSpace
     );
     LocalMemArrayGV->setUnnamedAddr(GlobalVariable::UnnamedAddr::Local);
     constexpr int DefaultMaxAlignment = 128;
-    if (!TT.isSPIROrSPIRV())
+    if (!usesKernelArgForDynWGLocalMem(TT))
       LocalMemArrayGV->setAlignment(Align{DefaultMaxAlignment});
   }
   lowerLocalMemCall(DLMFunc, [&](CallInst *CI) {
     lowerDynamicLocalMemCallDirect(CI, TT, LocalMemArrayGV);
   });
-  if (TT.isSPIROrSPIRV()) {
+  if (usesKernelArgForDynWGLocalMem(TT)) {
     SmallVector<Function *, 4> Kernels;
     llvm::for_each(M.functions(), [&](Function &F) {
       if (F.getCallingConv() == CallingConv::SPIR_KERNEL &&