[SYCL] Allow specification of double GRF mode for SYCL (#6914)

This change extends Konst's work from
#6182 to work for any SYCL kernel, not
just ESIMD kernels

Basic summary of changes:
1) Move SYCL library set_kernel_properties function and related detail
header out of esimd code into generic SYCL code
2) Generalize SYCLLowerESIMDKernelPropsPass to make it work for SYCL
kernels
3) Change sycl-post-link module splitting to split non-ESIMD modules
that have any number of double GRF kernels
4) Change program loader to add the "-ze-opt-large-register-file" option
if the double GRF property is set. We do this instead of -doubleGRF
because -doubleGRF only works for the VC backend, while
-ze-opt-large-register-file works for both VC and scalar backends

Signed-off-by: Sarnie, Nick <nick.sarnie@intel.com>

llvm/include/llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h

@@ -11,43 +11,11 @@
 #include "llvm/ADT/SmallPtrSet.h"
 #include "llvm/IR/Function.h"
 
-#include <functional>
-
 namespace llvm {
 namespace esimd {
 
-constexpr char ATTR_DOUBLE_GRF[] = "esimd-double-grf";
 constexpr char ESIMD_MARKER_MD[] = "sycl_explicit_simd";
 
-using CallGraphNodeAction = std::function<void(Function *)>;
-
-// Traverses call graph starting from given function up the call chain applying
-// given action to each function met on the way. If \c ErrorOnNonCallUse
-// parameter is true, then no functions' uses are allowed except calls.
-// Otherwise, any function where use of the current one happened is added to the
-// call graph as if the use was a call.
-// Functions which are part of the visited set ('Visited' parameter) are not
-// traversed.
-void traverseCallgraphUp(llvm::Function *F, CallGraphNodeAction NodeF,
-                         SmallPtrSetImpl<Function *> &Visited,
-                         bool ErrorOnNonCallUse);
-
-template <class CallGraphNodeActionF>
-void traverseCallgraphUp(Function *F, CallGraphNodeActionF ActionF,
-                         SmallPtrSetImpl<Function *> &Visited,
-                         bool ErrorOnNonCallUse) {
-  traverseCallgraphUp(F, CallGraphNodeAction(ActionF), Visited,
-                      ErrorOnNonCallUse);
-}
-
-template <class CallGraphNodeActionF>
-void traverseCallgraphUp(Function *F, CallGraphNodeActionF ActionF,
-                         bool ErrorOnNonCallUse = true) {
-  SmallPtrSet<Function *, 32> Visited;
-  traverseCallgraphUp(F, CallGraphNodeAction(ActionF), Visited,
-                      ErrorOnNonCallUse);
-}
-
 // Tells whether given function is a ESIMD kernel.
 bool isESIMDKernel(const Function &F);
 // Tells whether given function is a ESIMD function.

llvm/include/llvm/SYCLLowerIR/LowerKernelProps.h

@@ -0,0 +1,27 @@
+//===---- LowerKernelProps.h - lower kernel properties -----------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+// Lowers SYCL kernel properties into attributes used by sycl-post-link
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include "llvm/IR/PassManager.h"
+
+namespace sycl {
+namespace kernel_props {
+constexpr char ATTR_DOUBLE_GRF[] = "double-grf";
+}
+} // namespace sycl
+namespace llvm {
+// Lowers calls to __sycl_set_kernel_properties
+class SYCLLowerKernelPropsPass
+    : public PassInfoMixin<SYCLLowerKernelPropsPass> {
+public:
+  PreservedAnalyses run(Module &M, ModuleAnalysisManager &);
+};
+} // namespace llvm

llvm/include/llvm/SYCLLowerIR/SYCLUtils.h

@@ -0,0 +1,50 @@
+//===------------ SYCLUtils.h - SYCL utility functions
+//------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+// Utility functions for SYCL.
+//===----------------------------------------------------------------------===//
+#pragma once
+
+#include "llvm/ADT/SmallPtrSet.h"
+#include "llvm/IR/Function.h"
+
+#include <functional>
+namespace llvm {
+namespace sycl {
+namespace utils {
+using CallGraphNodeAction = std::function<void(Function *)>;
+
+// Traverses call graph starting from given function up the call chain applying
+// given action to each function met on the way. If \c ErrorOnNonCallUse
+// parameter is true, then no functions' uses are allowed except calls.
+// Otherwise, any function where use of the current one happened is added to the
+// call graph as if the use was a call.
+// Functions which are part of the visited set ('Visited' parameter) are not
+// traversed.
+void traverseCallgraphUp(llvm::Function *F, CallGraphNodeAction NodeF,
+                         SmallPtrSetImpl<Function *> &Visited,
+                         bool ErrorOnNonCallUse);
+
+template <class CallGraphNodeActionF>
+void traverseCallgraphUp(Function *F, CallGraphNodeActionF ActionF,
+                         SmallPtrSetImpl<Function *> &Visited,
+                         bool ErrorOnNonCallUse) {
+  traverseCallgraphUp(F, CallGraphNodeAction(ActionF), Visited,
+                      ErrorOnNonCallUse);
+}
+
+template <class CallGraphNodeActionF>
+void traverseCallgraphUp(Function *F, CallGraphNodeActionF ActionF,
+                         bool ErrorOnNonCallUse = true) {
+  SmallPtrSet<Function *, 32> Visited;
+  traverseCallgraphUp(F, CallGraphNodeAction(ActionF), Visited,
+                      ErrorOnNonCallUse);
+}
+} // namespace utils
+} // namespace sycl
+} // namespace llvm

llvm/lib/Passes/PassBuilder.cpp

@@ -81,6 +81,7 @@
 #include "llvm/SYCLLowerIR/ESIMD/ESIMDVerifier.h"
 #include "llvm/SYCLLowerIR/ESIMD/LowerESIMD.h"
 #include "llvm/SYCLLowerIR/LowerInvokeSimd.h"
+#include "llvm/SYCLLowerIR/LowerKernelProps.h"
 #include "llvm/SYCLLowerIR/LowerWGLocalMemory.h"
 #include "llvm/SYCLLowerIR/LowerWGScope.h"
 #include "llvm/SYCLLowerIR/MutatePrintfAddrspace.h"

llvm/lib/Passes/PassRegistry.def

@@ -127,7 +127,7 @@ MODULE_PASS("memprof-module", ModuleMemProfilerPass())
 MODULE_PASS("poison-checking", PoisonCheckingPass())
 MODULE_PASS("pseudo-probe-update", PseudoProbeUpdatePass())
 MODULE_PASS("LowerESIMD", SYCLLowerESIMDPass())
-MODULE_PASS("lower-esimd-kernel-props", SYCLLowerESIMDKernelPropsPass())
+MODULE_PASS("lower-kernel-props", SYCLLowerKernelPropsPass())
 MODULE_PASS("ESIMDLowerVecArg", ESIMDLowerVecArgPass())
 MODULE_PASS("esimd-opt-call-conv", ESIMDOptimizeVecArgCallConvPass())
 MODULE_PASS("esimd-verifier", ESIMDVerifierPass())

llvm/lib/SYCLLowerIR/CMakeLists.txt

@@ -48,18 +48,19 @@ set_property(GLOBAL PROPERTY LLVMGenXIntrinsics_BINARY_PROP ${LLVMGenXIntrinsics
 
 add_llvm_component_library(LLVMSYCLLowerIR
   ESIMD/LowerESIMD.cpp
-  ESIMD/LowerESIMDKernelProps.cpp
   ESIMD/LowerESIMDVLoadVStore.cpp
   ESIMD/LowerESIMDVecArg.cpp
   ESIMD/ESIMDUtils.cpp
   ESIMD/ESIMDVerifier.cpp
   ESIMD/LowerESIMDKernelAttrs.cpp
   ESIMD/ESIMDOptimizeVecArgCallConv.cpp
   LowerInvokeSimd.cpp
+  LowerKernelProps.cpp
   LowerWGScope.cpp
   LowerWGLocalMemory.cpp
   MutatePrintfAddrspace.cpp
   SYCLPropagateAspectsUsage.cpp
+  SYCLUtils.cpp
 
   LocalAccessorToSharedMemory.cpp
   GlobalOffset.cpp

llvm/lib/SYCLLowerIR/ESIMD/ESIMDUtils.cpp

@@ -19,61 +19,6 @@
 namespace llvm {
 namespace esimd {
 
-void traverseCallgraphUp(llvm::Function *F, CallGraphNodeAction ActionF,
-                         SmallPtrSetImpl<Function *> &FunctionsVisited,
-                         bool ErrorOnNonCallUse) {
-  SmallVector<Function *, 32> Worklist;
-
-  if (FunctionsVisited.count(F) == 0)
-    Worklist.push_back(F);
-
-  while (!Worklist.empty()) {
-    Function *CurF = Worklist.pop_back_val();
-    FunctionsVisited.insert(CurF);
-    // Apply the action function.
-    ActionF(CurF);
-
-    // Update all callers as well.
-    for (auto It = CurF->use_begin(); It != CurF->use_end(); It++) {
-      auto FCall = It->getUser();
-      auto ErrMsg =
-          llvm::Twine(__FILE__ " ") +
-          "Function use other than call detected while traversing call\n"
-          "graph up to a kernel";
-      if (!isa<CallInst>(FCall)) {
-        // A use other than a call is met...
-        if (ErrorOnNonCallUse) {
-          // ... non-call is an error - report
-          llvm::report_fatal_error(ErrMsg);
-        } else {
-          // ... non-call is OK - add using function to the worklist
-          if (auto *I = dyn_cast<Instruction>(FCall)) {
-            auto UseF = I->getFunction();
-
-            if (FunctionsVisited.count(UseF) == 0) {
-              Worklist.push_back(UseF);
-            }
-          }
-        }
-      } else {
-        auto *CI = cast<CallInst>(FCall);
-
-        if ((CI->getCalledFunction() != CurF)) {
-          // CurF is used in a call, but not as the callee.
-          if (ErrorOnNonCallUse)
-            llvm::report_fatal_error(ErrMsg);
-        } else {
-          auto FCaller = CI->getFunction();
-
-          if (!FunctionsVisited.count(FCaller)) {
-            Worklist.push_back(FCaller);
-          }
-        }
-      }
-    }
-  }
-}
-
 bool isESIMD(const Function &F) {
   return F.getMetadata(ESIMD_MARKER_MD) != nullptr;
 }

llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp

@@ -48,6 +48,7 @@ static const char *LegalSYCLFunctions[] = {
     "^sycl::_V1::exp<.+>",
     "^sycl::_V1::bit_cast<.+>",
     "^sycl::_V1::operator.+<.+>",
+    "^sycl::_V1::ext::intel::experimental::set_kernel_properties",
     "^sycl::_V1::ext::oneapi::sub_group::.+",
     "^sycl::_V1::ext::oneapi::experimental::spec_constant<.+>::.+",
     "^sycl::_V1::ext::oneapi::experimental::this_sub_group",

llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp

@@ -15,6 +15,7 @@
 
 #include "llvm/SYCLLowerIR/ESIMD/LowerESIMD.h"
 #include "llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h"
+#include "llvm/SYCLLowerIR/SYCLUtils.h"
 
 #include "llvm/ADT/DenseMap.h"
 #include "llvm/ADT/DenseSet.h"
@@ -977,7 +978,7 @@ static void translateSLMInit(CallInst &CI) {
       *F->getParent(), genx::KernelMDOp::SLMSize, NewVal};
   // TODO: Keep track of traversed functions (use 4-argument version of
   // traverseCallgraphUp) to avoid repeating traversals over same function.
-  esimd::traverseCallgraphUp(F, SetMaxSLMSize);
+  sycl::utils::traverseCallgraphUp(F, SetMaxSLMSize);
 }
 
 // This function sets/updates VCNamedBarrierCount attribute to the kernels
@@ -995,7 +996,7 @@ static void translateNbarrierInit(CallInst &CI) {
       *F->getParent(), genx::KernelMDOp::NBarrierCnt, NewVal};
   // TODO: Keep track of traversed functions to avoid repeating traversals
   // over same function.
-  esimd::traverseCallgraphUp(F, SetMaxNBarrierCnt);
+  sycl::utils::traverseCallgraphUp(F, SetMaxNBarrierCnt);
 }
 
 static void translatePackMask(CallInst &CI) {
@@ -1771,8 +1772,8 @@ size_t SYCLLowerESIMDPass::runOnFunction(Function &F,
         ToErase.push_back(CI);
         continue;
       }
-      assert(!Name.startswith("__esimd_set_kernel_properties") &&
-             "__esimd_set_kernel_properties must have been lowered");
+      assert(!Name.startswith("__sycl_set_kernel_properties") &&
+             "__sycl_set_kernel_properties must have been lowered");
 
       if (Name.empty() || !Name.startswith(ESIMD_INTRIN_PREF1))
         continue;

llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp

@@ -10,6 +10,7 @@
 
 #include "llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h"
 #include "llvm/SYCLLowerIR/ESIMD/LowerESIMD.h"
+#include "llvm/SYCLLowerIR/SYCLUtils.h"
 
 #include "llvm/IR/Module.h"
 #include "llvm/Pass.h"
@@ -26,7 +27,7 @@ SYCLFixupESIMDKernelWrapperMDPass::run(Module &M, ModuleAnalysisManager &MAM) {
     if (llvm::esimd::isESIMD(F)) {
       // TODO: Keep track of traversed functions to avoid repeating traversals
       // over same function.
-      llvm::esimd::traverseCallgraphUp(
+      sycl::utils::traverseCallgraphUp(
           &F,
           [&](Function *GraphNode) {
             if (!llvm::esimd::isESIMD(*GraphNode)) {

llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelProps.cpp → llvm/lib/SYCLLowerIR/LowerKernelProps.cpp

@@ -1,34 +1,34 @@
-//===---- LowerESIMDKernelProps.h - lower __esimd_set_kernel_properties ---===//
+//===---- LowerKernelProps.cpp - lower __sycl_set_kernel_properties ---===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
-// Finds and lowers __esimd_set_kernel_properties calls: converts the call to
+// Finds and lowers __sycl_set_kernel_properties calls: converts the call to
 // function attributes and adds those attributes to all kernels which can
 // potentially call this intrinsic.
 
-#include "llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h"
-#include "llvm/SYCLLowerIR/ESIMD/LowerESIMD.h"
+#include "llvm/SYCLLowerIR/LowerKernelProps.h"
+#include "llvm/SYCLLowerIR/SYCLUtils.h"
 
 #include "llvm/ADT/SmallPtrSet.h"
 #include "llvm/IR/Instructions.h"
 #include "llvm/IR/Module.h"
 #include "llvm/IR/Operator.h"
 #include "llvm/Pass.h"
 
-#define DEBUG_TYPE "LowerESIMDKernelProps"
+#define DEBUG_TYPE "LowerKernelProps"
 
 using namespace llvm;
 
 namespace {
 
 constexpr char SET_KERNEL_PROPS_FUNC_NAME[] =
-    "_Z29__esimd_set_kernel_propertiesi";
+    "_Z28__sycl_set_kernel_propertiesi";
 
 // Kernel property identifiers. Should match ones in
-// sycl/include/sycl/ext/intel/experimental/esimd/kernel_properties.hpp
+// sycl/include/sycl/ext/intel/experimental/kernel_properties.hpp
 enum property_ids { use_double_grf = 0 };
 
 void processSetKernelPropertiesCall(CallInst &CI) {
@@ -46,10 +46,8 @@ void processSetKernelPropertiesCall(CallInst &CI) {
   case property_ids::use_double_grf:
     // TODO: Keep track of traversed functions to avoid repeating traversals
     // over same function.
-    llvm::esimd::traverseCallgraphUp(F, [](Function *GraphNode) {
-      if (llvm::esimd::isESIMDKernel(*GraphNode)) {
-        GraphNode->addFnAttr(llvm::esimd::ATTR_DOUBLE_GRF);
-      }
+    llvm::sycl::utils::traverseCallgraphUp(F, [](Function *GraphNode) {
+      GraphNode->addFnAttr(::sycl::kernel_props::ATTR_DOUBLE_GRF);
     });
     break;
   default:
@@ -60,8 +58,8 @@ void processSetKernelPropertiesCall(CallInst &CI) {
 } // namespace
 
 namespace llvm {
-PreservedAnalyses
-SYCLLowerESIMDKernelPropsPass::run(Module &M, ModuleAnalysisManager &MAM) {
+PreservedAnalyses SYCLLowerKernelPropsPass::run(Module &M,
+                                                ModuleAnalysisManager &MAM) {
   Function *F = M.getFunction(SET_KERNEL_PROPS_FUNC_NAME);
 
   if (!F) {
@@ -71,7 +69,7 @@ SYCLLowerESIMDKernelPropsPass::run(Module &M, ModuleAnalysisManager &MAM) {
   SmallVector<User *, 4> Users(F->users());
 
   for (User *Usr : Users) {
-    // a call can be the only use of the __esimd_set_kernel_properties built-in
+    // a call can be the only use of the __sycl_set_kernel_properties built-in
     CallInst *CI = cast<CallInst>(Usr);
     processSetKernelPropertiesCall(*CI);
     CI->eraseFromParent();