[SYCL] Allow specification of double GRF mode for SYCL

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

@@ -16,7 +16,6 @@
 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 *)>;

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,14 @@ 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

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

@@ -1771,8 +1771,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/LowerESIMDKernelProps.cpp → llvm/lib/SYCLLowerIR/LowerKernelProps.cpp

@@ -1,34 +1,34 @@
-//===---- LowerESIMDKernelProps.h - lower __esimd_set_kernel_properties ---===//
+//===---- LowerESIMDKernelProps.h - 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/LowerKernelProps.h"
 #include "llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h"
-#include "llvm/SYCLLowerIR/ESIMD/LowerESIMD.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) {
@@ -47,9 +47,7 @@ void processSetKernelPropertiesCall(CallInst &CI) {
     // 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);
-      }
+      GraphNode->addFnAttr(llvm::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();

llvm/test/SYCLLowerIR/lower_kernel_props.ll

@@ -0,0 +1,44 @@
+; This test checks handling of the
+;   __sycl_set_kernel_properties(...);
+; intrinsic by LowerKernelProps pass - it should:
+; - determine kernels calling this intrinsic (walk up the call graph)
+; - remove the intrinsic call 
+; - mark the kernel with corresponding attribute (only "double-grf" for now)
+
+; RUN: opt -passes=lower-kernel-props -S %s -o - | FileCheck %s
+
+; ModuleID = 'double_grf.bc'
+source_filename = "llvm-link"
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
+target triple = "spir64-unknown-unknown"
+
+define dso_local spir_func void @_Z17double_grf_markerv() {
+; CHECK: define dso_local spir_func void @_Z17double_grf_markerv()
+; -- '0' constant argument means "double GRF" property:
+  call spir_func void @_Z28__sycl_set_kernel_propertiesi(i32 noundef 0)
+; -- Check that LowerKernelProps removed the marker call above:
+; CHECK-NOT: {{.*}} @_Z28__sycl_set_kernel_propertiesi
+  ret void
+; CHECK-NEXT: ret void
+}
+
+declare dso_local spir_func void @_Z28__sycl_set_kernel_propertiesi(i32 noundef)
+
+; -- This kernel calls the marker function indirectly
+define weak_odr dso_local spir_kernel void @__double_grf_kernel1() !sycl_explicit_simd !0 !intel_reqd_sub_group_size !1 {
+; CHECK: {{.*}} spir_kernel void @__double_grf_kernel1() #0
+  call spir_func void @_Z17double_grf_markerv()
+  ret void
+}
+
+; -- This kernel calls the marker function directly
+define weak_odr dso_local spir_kernel void @__double_grf_kernel2() #0 !sycl_explicit_simd !0 !intel_reqd_sub_group_size !1 {
+; CHECK: {{.*}} spir_kernel void @__double_grf_kernel2() #0
+  call spir_func void @_Z28__sycl_set_kernel_propertiesi(i32 noundef 0)
+  ret void
+}
+
+attributes #0 = { "double-grf" }
+
+!0 = !{}
+!1 = !{i32 1}

llvm/test/tools/sycl-post-link/sycl-double-grf.ll

@@ -0,0 +1,58 @@
+; This test checks handling of the
+;   set_kernel_properties(kernel_properties::use_double_grf);
+; by the post-link-tool:
+; - ESIMD/SYCL splitting happens as usual
+; - ESIMD module is further split into callgraphs for entry points requesting
+;   "double GRF" and callgraphs for entry points which are not
+; - Compiler adds 'isDoubleGRF' property to the device binary
+;   images requesting "double GRF" 
+
+; RUN: sycl-post-link -split=source -symbols -split-esimd -lower-esimd -S %s -o %t.table
+; RUN: FileCheck %s -input-file=%t.table
+; RUN: FileCheck %s -input-file=%t_x2grf_0.ll --check-prefixes CHECK-2xGRF-IR
+; RUN: FileCheck %s -input-file=%t_x2grf_0.prop --check-prefixes CHECK-2xGRF-PROP
+; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-SYCL-SYM
+; RUN: FileCheck %s -input-file=%t_x2grf_0.sym --check-prefixes CHECK-2xGRF-SYM
+
+; CHECK: [Code|Properties|Symbols]
+; CHECK: {{.*}}_x2grf_0.ll|{{.*}}_x2grf_0.prop|{{.*}}_x2grf_0.sym
+; CHECK: {{.*}}_0.ll|{{.*}}_0.prop|{{.*}}_0.sym
+
+; CHECK-2xGRF-PROP: isDoubleGRF=1|1
+
+; CHECK-SYCL-SYM: __SYCL_kernel
+; CHECK-SYCL-SYM-EMPTY:
+
+; CHECK-2xGRF-SYM: __double_grf_kernel
+; CHECK-2xGRF-SYM-EMPTY:
+
+; ModuleID = 'double_grf.bc'
+source_filename = "llvm-link"
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
+target triple = "spir64-unknown-unknown"
+
+define weak_odr dso_local spir_kernel void @__SYCL_kernel() #0 {
+entry:
+  ret void
+}
+
+define dso_local spir_func void @_Z17double_grf_markerv() {
+entry:
+  call spir_func void @_Z28__sycl_set_kernel_propertiesi(i32 noundef 0)
+; -- Check that LowerKernelProps lowering removed the marker call above:
+; CHECK-2xGRF-IR-NOT: {{.*}} @_Z28__sycl_set_kernel_propertiesi
+  ret void
+}
+
+declare dso_local spir_func void @_Z28__sycl_set_kernel_propertiesi(i32 noundef)
+
+define weak_odr dso_local spir_kernel void @__double_grf_kernel() #0 {
+entry:
+  call spir_func void @_Z17double_grf_markerv()
+  ret void
+}
+
+attributes #0 = { "sycl-module-id"="a.cpp" }
+
+!0 = !{}
+!1 = !{i32 1}

llvm/test/tools/sycl-post-link/sycl-esimd/sycl-esimd-double-grf.ll → llvm/test/tools/sycl-post-link/sycl-esimd-double-grf.ll

@@ -4,7 +4,7 @@
 ; - ESIMD/SYCL splitting happens as usual
 ; - ESIMD module is further split into callgraphs for entry points requesting
 ;   "double GRF" and callgraphs for entry points which are not
-; - Compiler adds 'isDoubleGRFEsimdImage' property to the ESIMD device binary
+; - Compiler adds 'isDoubleGRF' property to the ESIMD device binary
 ;   images requesting "double GRF" 
 
 ; RUN: sycl-post-link -split=source -symbols -split-esimd -lower-esimd -S %s -o %t.table
@@ -21,7 +21,7 @@
 ; CHECK: {{.*}}esimd_0.ll|{{.*}}esimd_0.prop|{{.*}}esimd_0.sym
 
 ; CHECK-ESIMD-2xGRF-PROP: isEsimdImage=1|1
-; CHECK-ESIMD-2xGRF-PROP: isDoubleGRFEsimdImage=1|1
+; CHECK-ESIMD-2xGRF-PROP: isDoubleGRF=1|1
 
 ; CHECK-SYCL-SYM: __SYCL_kernel
 ; CHECK-SYCL-SYM-EMPTY:
@@ -49,13 +49,13 @@ entry:
 
 define dso_local spir_func void @_Z17double_grf_markerv() {
 entry:
-  call spir_func void @_Z29__esimd_set_kernel_propertiesi(i32 noundef 0)
+  call spir_func void @_Z28__sycl_set_kernel_propertiesi(i32 noundef 0)
 ; -- Check that ESIMD lowering removed the marker call above:
-; CHECK-ESIMD-2xGRF-IR-NOT: {{.*}} @_Z29__esimd_set_kernel_propertiesi
+; CHECK-ESIMD-2xGRF-IR-NOT: {{.*}} @_Z28__sycl_set_kernel_propertiesi
   ret void
 }
 
-declare dso_local spir_func void @_Z29__esimd_set_kernel_propertiesi(i32 noundef)
+declare dso_local spir_func void @_Z28__sycl_set_kernel_propertiesi(i32 noundef)
 
 define weak_odr dso_local spir_kernel void @__ESIMD_double_grf_kernel() #0 !sycl_explicit_simd !0 !intel_reqd_sub_group_size !1 {
 entry:

llvm/tools/sycl-post-link/ModuleSplitter.cpp

@@ -19,6 +19,7 @@
 #include "llvm/IR/LegacyPassManager.h"
 #include "llvm/IR/Module.h"
 #include "llvm/SYCLLowerIR/LowerInvokeSimd.h"
+#include "llvm/SYCLLowerIR/LowerKernelProps.h"
 #include "llvm/Transforms/IPO.h"
 #include "llvm/Transforms/IPO/GlobalDCE.h"
 #include "llvm/Transforms/IPO/StripDeadPrototypes.h"
@@ -41,7 +42,6 @@ constexpr char ESIMD_SCOPE_NAME[] = "<ESIMD>";
 constexpr char ESIMD_MARKER_MD[] = "sycl_explicit_simd";
 
 constexpr char ATTR_SYCL_MODULE_ID[] = "sycl-module-id";
-constexpr char ATTR_DOUBLE_GRF[] = "esimd-double-grf";
 
 bool hasIndirectFunctionsOrCalls(const Module &M) {
   for (const auto &F : M.functions()) {
@@ -726,11 +726,11 @@ void EntryPointGroup::rebuildFromNames(const std::vector<std::string> &Names,
 }
 
 std::unique_ptr<ModuleSplitterBase>
-getESIMDDoubleGRFSplitter(ModuleDesc &&MD, bool EmitOnlyKernelsAsEntryPoints) {
+getDoubleGRFSplitter(ModuleDesc &&MD, bool EmitOnlyKernelsAsEntryPoints) {
   EntryPointGroupVec Groups = groupEntryPointsByAttribute(
-      MD, ATTR_DOUBLE_GRF, EmitOnlyKernelsAsEntryPoints,
-      [](EntryPointGroup &G) {
-        if (G.GroupId == ATTR_DOUBLE_GRF) {
+      MD, llvm::sycl_kernel_props::ATTR_DOUBLE_GRF,
+      EmitOnlyKernelsAsEntryPoints, [](EntryPointGroup &G) {
+        if (G.GroupId == llvm::sycl_kernel_props::ATTR_DOUBLE_GRF) {
           G.Props.UsesDoubleGRF = true;
         }
       });

llvm/tools/sycl-post-link/ModuleSplitter.h

@@ -251,7 +251,7 @@ getSplitterByMode(ModuleDesc &&MD, IRSplitMode Mode,
                   bool EmitOnlyKernelsAsEntryPoints);
 
 std::unique_ptr<ModuleSplitterBase>
-getESIMDDoubleGRFSplitter(ModuleDesc &&MD, bool EmitOnlyKernelsAsEntryPoints);
+getDoubleGRFSplitter(ModuleDesc &&MD, bool EmitOnlyKernelsAsEntryPoints);
 
 #ifndef NDEBUG
 void dumpEntryPoints(const EntryPointSet &C, const char *msg = "", int Tab = 0);