diff --git a/llvm/include/llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h b/llvm/include/llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h index e07e61f34ed87..3aa6f8eca23bb 100644 --- a/llvm/include/llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h +++ b/llvm/include/llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h @@ -11,43 +11,11 @@ #include "llvm/ADT/SmallPtrSet.h" #include "llvm/IR/Function.h" -#include - namespace llvm { namespace esimd { -constexpr char ATTR_DOUBLE_GRF[] = "esimd-double-grf"; constexpr char ESIMD_MARKER_MD[] = "sycl_explicit_simd"; -using CallGraphNodeAction = std::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 &Visited, - bool ErrorOnNonCallUse); - -template -void traverseCallgraphUp(Function *F, CallGraphNodeActionF ActionF, - SmallPtrSetImpl &Visited, - bool ErrorOnNonCallUse) { - traverseCallgraphUp(F, CallGraphNodeAction(ActionF), Visited, - ErrorOnNonCallUse); -} - -template -void traverseCallgraphUp(Function *F, CallGraphNodeActionF ActionF, - bool ErrorOnNonCallUse = true) { - SmallPtrSet 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. diff --git a/llvm/include/llvm/SYCLLowerIR/LowerKernelProps.h b/llvm/include/llvm/SYCLLowerIR/LowerKernelProps.h new file mode 100644 index 0000000000000..81e12f9c93d9e --- /dev/null +++ b/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 { +public: + PreservedAnalyses run(Module &M, ModuleAnalysisManager &); +}; +} // namespace llvm diff --git a/llvm/include/llvm/SYCLLowerIR/SYCLUtils.h b/llvm/include/llvm/SYCLLowerIR/SYCLUtils.h new file mode 100644 index 0000000000000..65b089a6233a2 --- /dev/null +++ b/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 +namespace llvm { +namespace sycl { +namespace utils { +using CallGraphNodeAction = std::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 &Visited, + bool ErrorOnNonCallUse); + +template +void traverseCallgraphUp(Function *F, CallGraphNodeActionF ActionF, + SmallPtrSetImpl &Visited, + bool ErrorOnNonCallUse) { + traverseCallgraphUp(F, CallGraphNodeAction(ActionF), Visited, + ErrorOnNonCallUse); +} + +template +void traverseCallgraphUp(Function *F, CallGraphNodeActionF ActionF, + bool ErrorOnNonCallUse = true) { + SmallPtrSet Visited; + traverseCallgraphUp(F, CallGraphNodeAction(ActionF), Visited, + ErrorOnNonCallUse); +} +} // namespace utils +} // namespace sycl +} // namespace llvm diff --git a/llvm/lib/Passes/PassBuilder.cpp b/llvm/lib/Passes/PassBuilder.cpp index dabaa3a04135c..a3ecb88881503 100644 --- a/llvm/lib/Passes/PassBuilder.cpp +++ b/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" diff --git a/llvm/lib/Passes/PassRegistry.def b/llvm/lib/Passes/PassRegistry.def index 040a691da109a..95ddca64d9db5 100644 --- a/llvm/lib/Passes/PassRegistry.def +++ b/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()) diff --git a/llvm/lib/SYCLLowerIR/CMakeLists.txt b/llvm/lib/SYCLLowerIR/CMakeLists.txt index ff65c2e9d25e3..b3b14afe77931 100644 --- a/llvm/lib/SYCLLowerIR/CMakeLists.txt +++ b/llvm/lib/SYCLLowerIR/CMakeLists.txt @@ -48,7 +48,6 @@ 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 @@ -56,10 +55,12 @@ add_llvm_component_library(LLVMSYCLLowerIR ESIMD/LowerESIMDKernelAttrs.cpp ESIMD/ESIMDOptimizeVecArgCallConv.cpp LowerInvokeSimd.cpp + LowerKernelProps.cpp LowerWGScope.cpp LowerWGLocalMemory.cpp MutatePrintfAddrspace.cpp SYCLPropagateAspectsUsage.cpp + SYCLUtils.cpp LocalAccessorToSharedMemory.cpp GlobalOffset.cpp diff --git a/llvm/lib/SYCLLowerIR/ESIMD/ESIMDUtils.cpp b/llvm/lib/SYCLLowerIR/ESIMD/ESIMDUtils.cpp index 7728ae328686c..1e26507ad0845 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/ESIMDUtils.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/ESIMDUtils.cpp @@ -19,61 +19,6 @@ namespace llvm { namespace esimd { -void traverseCallgraphUp(llvm::Function *F, CallGraphNodeAction ActionF, - SmallPtrSetImpl &FunctionsVisited, - bool ErrorOnNonCallUse) { - SmallVector 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(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(FCall)) { - auto UseF = I->getFunction(); - - if (FunctionsVisited.count(UseF) == 0) { - Worklist.push_back(UseF); - } - } - } - } else { - auto *CI = cast(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; } diff --git a/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp b/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp index 766373e330834..ce82df3760bae 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp +++ b/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", diff --git a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp index ea35c02a7530d..bcaa462a3dde5 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp +++ b/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; diff --git a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp index f74544133e82e..ccf7da4ed8a53 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp +++ b/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)) { diff --git a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelProps.cpp b/llvm/lib/SYCLLowerIR/LowerKernelProps.cpp similarity index 70% rename from llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelProps.cpp rename to llvm/lib/SYCLLowerIR/LowerKernelProps.cpp index 684f7c47ebc06..2eb227b29e34c 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelProps.cpp +++ b/llvm/lib/SYCLLowerIR/LowerKernelProps.cpp @@ -1,16 +1,16 @@ -//===---- 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" @@ -18,17 +18,17 @@ #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 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(Usr); processSetKernelPropertiesCall(*CI); CI->eraseFromParent(); diff --git a/llvm/lib/SYCLLowerIR/SYCLUtils.cpp b/llvm/lib/SYCLLowerIR/SYCLUtils.cpp new file mode 100644 index 0000000000000..5eaf6a9a02131 --- /dev/null +++ b/llvm/lib/SYCLLowerIR/SYCLUtils.cpp @@ -0,0 +1,73 @@ +//===------------ SYCLUtils.cpp - 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. +//===----------------------------------------------------------------------===// +#include "llvm/SYCLLowerIR/SYCLUtils.h" +#include "llvm/IR/Instructions.h" + +namespace llvm { +namespace sycl { +namespace utils { +void traverseCallgraphUp(llvm::Function *F, CallGraphNodeAction ActionF, + SmallPtrSetImpl &FunctionsVisited, + bool ErrorOnNonCallUse) { + SmallVector 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(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(FCall)) { + auto UseF = I->getFunction(); + + if (FunctionsVisited.count(UseF) == 0) { + Worklist.push_back(UseF); + } + } + } + } else { + auto *CI = cast(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 utils +} // namespace sycl +} // namespace llvm diff --git a/llvm/test/SYCLLowerIR/ESIMD/lower_kernel_props.ll b/llvm/test/SYCLLowerIR/ESIMD/lower_kernel_props.ll deleted file mode 100644 index d9b2fe06cde20..0000000000000 --- a/llvm/test/SYCLLowerIR/ESIMD/lower_kernel_props.ll +++ /dev/null @@ -1,44 +0,0 @@ -; This test checks handling of the -; __esimd_set_kernel_properties(...); -; intrinsic by LowerESIMDKernelProps pass - it should: -; - determine kernels calling this intrinsic (walk up the call graph) -; - remove the intrinsic call -; - mark the kernel with corresponding attribute (only "esimd-double-grf" for now) - -; RUN: opt -passes=lower-esimd-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 @_Z29__esimd_set_kernel_propertiesi(i32 noundef 0) -; -- Check that ESIMD lowering removed the marker call above: -; CHECK-NOT: {{.*}} @_Z29__esimd_set_kernel_propertiesi - ret void -; CHECK-NEXT: ret void -} - -declare dso_local spir_func void @_Z29__esimd_set_kernel_propertiesi(i32 noundef) - -; -- This kernel calls the marker function indirectly -define weak_odr dso_local spir_kernel void @__ESIMD_double_grf_kernel1() !sycl_explicit_simd !0 !intel_reqd_sub_group_size !1 { -; CHECK: {{.*}} spir_kernel void @__ESIMD_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 @__ESIMD_double_grf_kernel2() #0 !sycl_explicit_simd !0 !intel_reqd_sub_group_size !1 { -; CHECK: {{.*}} spir_kernel void @__ESIMD_double_grf_kernel2() #0 - call spir_func void @_Z29__esimd_set_kernel_propertiesi(i32 noundef 0) - ret void -} - -attributes #0 = { "esimd-double-grf" } - -!0 = !{} -!1 = !{i32 1} diff --git a/llvm/test/SYCLLowerIR/lower_kernel_props.ll b/llvm/test/SYCLLowerIR/lower_kernel_props.ll new file mode 100644 index 0000000000000..aa36cd8c94c06 --- /dev/null +++ b/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} diff --git a/llvm/test/tools/sycl-post-link/sycl-double-grf.ll b/llvm/test/tools/sycl-post-link/sycl-double-grf.ll new file mode 100644 index 0000000000000..dcc5695ce48e9 --- /dev/null +++ b/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} diff --git a/llvm/test/tools/sycl-post-link/sycl-esimd/sycl-esimd-double-grf.ll b/llvm/test/tools/sycl-post-link/sycl-esimd-double-grf.ll similarity index 86% rename from llvm/test/tools/sycl-post-link/sycl-esimd/sycl-esimd-double-grf.ll rename to llvm/test/tools/sycl-post-link/sycl-esimd-double-grf.ll index 0abc1150b9634..5d6104a40505f 100644 --- a/llvm/test/tools/sycl-post-link/sycl-esimd/sycl-esimd-double-grf.ll +++ b/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: diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.cpp b/llvm/tools/sycl-post-link/ModuleSplitter.cpp index 0670c80efea94..a4ed14841a196 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.cpp +++ b/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[] = ""; 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()) { @@ -744,11 +744,11 @@ void EntryPointGroup::rebuildFromNames(const std::vector &Names, } std::unique_ptr -getESIMDDoubleGRFSplitter(ModuleDesc &&MD, bool EmitOnlyKernelsAsEntryPoints) { +getDoubleGRFSplitter(ModuleDesc &&MD, bool EmitOnlyKernelsAsEntryPoints) { EntryPointGroupVec Groups = groupEntryPointsByAttribute( - MD, ATTR_DOUBLE_GRF, EmitOnlyKernelsAsEntryPoints, + MD, sycl::kernel_props::ATTR_DOUBLE_GRF, EmitOnlyKernelsAsEntryPoints, [](EntryPointGroup &G) { - if (G.GroupId == ATTR_DOUBLE_GRF) { + if (G.GroupId == sycl::kernel_props::ATTR_DOUBLE_GRF) { G.Props.UsesDoubleGRF = true; } }); diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.h b/llvm/tools/sycl-post-link/ModuleSplitter.h index b5847b53b2777..f362c2c1973da 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.h +++ b/llvm/tools/sycl-post-link/ModuleSplitter.h @@ -251,7 +251,7 @@ getSplitterByMode(ModuleDesc &&MD, IRSplitMode Mode, bool EmitOnlyKernelsAsEntryPoints); std::unique_ptr -getESIMDDoubleGRFSplitter(ModuleDesc &&MD, bool EmitOnlyKernelsAsEntryPoints); +getDoubleGRFSplitter(ModuleDesc &&MD, bool EmitOnlyKernelsAsEntryPoints); #ifndef NDEBUG void dumpEntryPoints(const EntryPointSet &C, const char *msg = "", int Tab = 0); diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index 0d84a42b4cb05..c99a0c398b1e7 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -39,6 +39,7 @@ #include "llvm/Passes/PassBuilder.h" #include "llvm/SYCLLowerIR/ESIMD/LowerESIMD.h" #include "llvm/SYCLLowerIR/LowerInvokeSimd.h" +#include "llvm/SYCLLowerIR/LowerKernelProps.h" #include "llvm/Support/CommandLine.h" #include "llvm/Support/FileSystem.h" #include "llvm/Support/InitLLVM.h" @@ -442,8 +443,7 @@ std::string saveModuleProperties(module_split::ModuleDesc &MD, PropSet[PropSetRegTy::SYCL_MISC_PROP].insert({"isEsimdImage", true}); } if (MD.isDoubleGRF()) - PropSet[PropSetRegTy::SYCL_MISC_PROP].insert( - {"isDoubleGRFEsimdImage", true}); + PropSet[PropSetRegTy::SYCL_MISC_PROP].insert({"isDoubleGRF", true}); { std::vector FuncNames = getKernelNamesUsingAssert(M); for (const StringRef &FName : FuncNames) @@ -558,6 +558,14 @@ bool lowerEsimdConstructs(module_split::ModuleDesc &MD) { return !Res.areAllPreserved(); } +// Compute the filename suffix for the module +StringRef getModuleSuffix(const module_split::ModuleDesc &MD) { + if (MD.isDoubleGRF()) { + return MD.isESIMD() ? "_esimd_x2grf" : "_x2grf"; + } + return MD.isESIMD() ? "_esimd" : ""; +} + // @param MD Module descriptor to save // @param IRFilename filename of already available IR component. If not empty, // IR component saving is skipped, and this file name is recorded as such in @@ -567,8 +575,7 @@ bool lowerEsimdConstructs(module_split::ModuleDesc &MD) { IrPropSymFilenameTriple saveModule(module_split::ModuleDesc &MD, int I, StringRef IRFilename = "") { IrPropSymFilenameTriple Res; - StringRef Suffix = - MD.isDoubleGRF() ? "_esimd_x2grf" : (MD.isESIMD() ? "_esimd" : ""); + StringRef Suffix = getModuleSuffix(MD); if (!IRFilename.empty()) { // don't save IR, just record the filename @@ -731,7 +738,7 @@ processInputModule(std::unique_ptr M) { // Lower kernel properties setting APIs before "double GRF" splitting, as: // - the latter uses the result of the former // - saves processing time - Modified |= runModulePass(*M); + Modified |= runModulePass(*M); DUMP_ENTRY_POINTS(*M, EmitOnlyKernelsAsEntryPoints, "Input"); @@ -768,8 +775,8 @@ processInputModule(std::unique_ptr M) { DUMP_ENTRY_POINTS(MDesc.entries(), MDesc.Name.c_str(), 1); std::unique_ptr DoubleGRFSplitter = - module_split::getESIMDDoubleGRFSplitter(std::move(MDesc), - EmitOnlyKernelsAsEntryPoints); + module_split::getDoubleGRFSplitter(std::move(MDesc), + EmitOnlyKernelsAsEntryPoints); const bool SplitByDoubleGRF = DoubleGRFSplitter->totalSplits() > 1; Modified |= SplitByDoubleGRF; diff --git a/sycl/include/sycl/ext/intel/esimd.hpp b/sycl/include/sycl/ext/intel/esimd.hpp index 36bd7f7662d7f..58758c109178a 100644 --- a/sycl/include/sycl/ext/intel/esimd.hpp +++ b/sycl/include/sycl/ext/intel/esimd.hpp @@ -87,7 +87,6 @@ #include #include #include -#include #include #include diff --git a/sycl/include/sycl/ext/intel/experimental/esimd/detail/misc_intrin.hpp b/sycl/include/sycl/ext/intel/experimental/detail/misc_intrin.hpp similarity index 59% rename from sycl/include/sycl/ext/intel/experimental/esimd/detail/misc_intrin.hpp rename to sycl/include/sycl/ext/intel/experimental/detail/misc_intrin.hpp index 608005c425b86..fbae9267158ed 100644 --- a/sycl/include/sycl/ext/intel/experimental/esimd/detail/misc_intrin.hpp +++ b/sycl/include/sycl/ext/intel/experimental/detail/misc_intrin.hpp @@ -1,18 +1,24 @@ -//==------------ misc_intrin.hpp - DPC++ Explicit SIMD API -----------------==// +//==------------ misc_intrin.hpp - SYCL 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 // //===----------------------------------------------------------------------===// -// Declares miscellaneous Explicit SIMD intrinsics. +// Declares miscellaneous SYCL intrinsics. //===----------------------------------------------------------------------===// #pragma once -/// @cond ESIMD_DETAIL +/// @cond SYCL_DETAIL -__ESIMD_INTRIN void __esimd_set_kernel_properties(int prop_mask) +#ifdef __SYCL_DEVICE_ONLY__ +#define __SYCL_INTRIN SYCL_EXTERNAL +#else +#define __SYCL_INTRIN inline +#endif // __SYCL_DEVICE_ONLY__ + +__SYCL_INTRIN void __sycl_set_kernel_properties(int prop_mask) #ifdef __SYCL_DEVICE_ONLY__ ; #else @@ -20,4 +26,4 @@ __ESIMD_INTRIN void __esimd_set_kernel_properties(int prop_mask) } // Only "double GRF" property is supported for now, safe to ignore on host. #endif // __SYCL_DEVICE_ONLY__ -/// @endcond ESIMD_DETAIL +/// @endcond SYCL_DETAIL diff --git a/sycl/include/sycl/ext/intel/experimental/esimd/kernel_properties.hpp b/sycl/include/sycl/ext/intel/experimental/kernel_properties.hpp similarity index 83% rename from sycl/include/sycl/ext/intel/experimental/esimd/kernel_properties.hpp rename to sycl/include/sycl/ext/intel/experimental/kernel_properties.hpp index 445a7ec4e70ac..786a396e921a6 100644 --- a/sycl/include/sycl/ext/intel/experimental/esimd/kernel_properties.hpp +++ b/sycl/include/sycl/ext/intel/experimental/kernel_properties.hpp @@ -1,4 +1,4 @@ -//==---------------- kernel_properties.hpp - DPC++ Explicit SIMD API -------==// +//==---------------- kernel_properties.hpp - SYCL 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. @@ -10,14 +10,13 @@ #pragma once -#include -#include +#include #include namespace sycl { __SYCL_INLINE_VER_NAMESPACE(_V1) { -namespace ext::intel::experimental::esimd { +namespace ext::intel::experimental { namespace kernel_properties { @@ -26,7 +25,7 @@ namespace kernel_properties { /// // Implementation note: ::value fields should match property IDs -// specified in llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp +// specified in llvm/lib/SYCLLowerIR/LowerKernelProps.cpp namespace detail { // Proxy to access private property classes' fields from the API code. @@ -65,16 +64,16 @@ void set_kernel_properties(KernelProps... props) { constexpr bool IsDoubleGRF = std::is_same_v; if constexpr (IsDoubleGRF) { - __esimd_set_kernel_properties( + __sycl_set_kernel_properties( kernel_properties::detail::proxy< kernel_properties::use_double_grf_tag>::value); } else { static_assert(IsDoubleGRF && - "set_kernel_properties: invalid ESIMD kernel property"); + "set_kernel_properties: invalid kernel property"); } }); } -} // namespace ext::intel::experimental::esimd +} // namespace ext::intel::experimental } // __SYCL_INLINE_VER_NAMESPACE(_V1) } // namespace sycl diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 5488c9692cb1c..7871f0a4e053e 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -404,10 +404,7 @@ static void appendCompileOptionsFromImage(std::string &CompileOpts, CompileOpts += std::string(TemporaryStr); } bool isEsimdImage = getUint32PropAsBool(Img, "isEsimdImage"); - bool isDoubleGRFEsimdImage = - getUint32PropAsBool(Img, "isDoubleGRFEsimdImage"); - assert((!isDoubleGRFEsimdImage || isEsimdImage) && - "doubleGRF applies only to ESIMD binary images"); + bool isDoubleGRF = getUint32PropAsBool(Img, "isDoubleGRF"); // The -vc-codegen option is always preserved for ESIMD kernels, regardless // of the contents SYCL_PROGRAM_COMPILE_OPTIONS environment variable. if (isEsimdImage) { @@ -419,9 +416,10 @@ static void appendCompileOptionsFromImage(std::string &CompileOpts, if (detail::SYCLConfig::get() == 0) CompileOpts += " -disable-finalizer-msg"; } - if (isDoubleGRFEsimdImage) { - assert(!CompileOpts.empty()); // -vc-codegen must be present - CompileOpts += " -doubleGRF"; + if (isDoubleGRF) { + if (!CompileOpts.empty()) + CompileOpts += " "; + CompileOpts += "-ze-opt-large-register-file"; } }