[SYCL] Allow specification of double GRF mode for SYCL

llvm/include/llvm/SYCLLowerIR/CallgraphUtils.h

@@ -0,0 +1,48 @@
+//===------------ CallgraphUtils.h - Callgraph 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 traversing callgraphs.
+//===----------------------------------------------------------------------===//
+#pragma once
+
+#include "llvm/ADT/SmallPtrSet.h"
+#include "llvm/IR/Function.h"
+
+#include <functional>
+namespace llvm {
+namespace CallgraphUtils {
+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 CallgraphUtils
+} // namespace llvm

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,28 @@
+//===---- 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 llvm {
+
+namespace sycl_kernel_props {
+constexpr char ATTR_DOUBLE_GRF[] = "double-grf";
+}
+
+// Lowers calls to __sycl_set_kernel_properties
+class SYCLLowerKernelPropsPass
+    : public PassInfoMixin<SYCLLowerKernelPropsPass> {
+public:
+  PreservedAnalyses run(Module &M, ModuleAnalysisManager &);
+};
+
+} // 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,14 +48,15 @@ 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
+  CallgraphUtils.cpp
   LowerInvokeSimd.cpp
+  LowerKernelProps.cpp
   LowerWGScope.cpp
   LowerWGLocalMemory.cpp
   MutatePrintfAddrspace.cpp

llvm/lib/SYCLLowerIR/CallgraphUtils.cpp

@@ -0,0 +1,71 @@
+//===------------ CallgraphUtils.cpp - Callgraph 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 traversing callgraphs.
+//===----------------------------------------------------------------------===//
+#include "llvm/SYCLLowerIR/CallgraphUtils.h"
+#include "llvm/IR/Instructions.h"
+
+namespace llvm {
+namespace CallgraphUtils {
+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);
+          }
+        }
+      }
+    }
+  }
+}
+} // namespace CallgraphUtils
+} // namespace llvm

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

@@ -14,6 +14,7 @@
 //===----------------------------------------------------------------------===//
 
 #include "llvm/SYCLLowerIR/ESIMD/LowerESIMD.h"
+#include "llvm/SYCLLowerIR/CallgraphUtils.h"
 #include "llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h"
 
 #include "llvm/ADT/DenseMap.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);
+  CallgraphUtils::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);
+  CallgraphUtils::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

@@ -8,6 +8,7 @@
 // Finds and adds  sycl_explicit_simd attributes to wrapper functions that wrap
 // ESIMD kernel functions
 
+#include "llvm/SYCLLowerIR/CallgraphUtils.h"
 #include "llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h"
 #include "llvm/SYCLLowerIR/ESIMD/LowerESIMD.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(
+      llvm::CallgraphUtils::traverseCallgraphUp(
           &F,
           [&](Function *GraphNode) {
             if (!llvm::esimd::isESIMD(*GraphNode)) {