From 382054c2e5caec5a2628dce1525dd02eba81093a Mon Sep 17 00:00:00 2001 From: "Sarnie, Nick" Date: Tue, 26 Nov 2024 14:27:52 -0800 Subject: [PATCH] [SYCL][ESIMD] Fix propagation of ESIMD attribute for inlined functions Signed-off-by: Sarnie, Nick --- .../llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h | 8 ++++++++ llvm/lib/SYCLLowerIR/ESIMD/ESIMDUtils.cpp | 10 ++++++++++ llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp | 13 ++++--------- .../ESIMD/LowerESIMDKernelAttrs.cpp | 18 ++++++++++++++++-- llvm/test/SYCLLowerIR/ESIMD/prop_metadata.ll | 17 +++++++++++++++++ 5 files changed, 55 insertions(+), 11 deletions(-) create mode 100644 llvm/test/SYCLLowerIR/ESIMD/prop_metadata.ll diff --git a/llvm/include/llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h b/llvm/include/llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h index 9066df0d96fd7..1ef7a90dc4846 100644 --- a/llvm/include/llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h +++ b/llvm/include/llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h @@ -25,6 +25,10 @@ constexpr char GENX_KERNEL_METADATA[] = "genx.kernels"; // sycl/ext/oneapi/experimental/invoke_simd.hpp::__builtin_invoke_simd // overloads instantiations: constexpr char INVOKE_SIMD_PREF[] = "_Z33__regcall3____builtin_invoke_simd"; +// The regexp for ESIMD intrinsics: +// /^_Z(\d+)__esimd_\w+/ +static constexpr char ESIMD_INTRIN_PREF0[] = "_Z"; +static constexpr char ESIMD_INTRIN_PREF1[] = "__esimd_"; bool isSlmAllocatorConstructor(const Function &F); bool isSlmAllocatorDestructor(const Function &F); @@ -133,5 +137,9 @@ struct UpdateUint64MetaDataToMaxValue { // functions has changed its attribute to alwaysinline. bool prepareForAlwaysInliner(Module &M); +// Remove manging from an ESIMD intrinsic function. +// Returns empty on pattern match failure. +StringRef stripMangling(StringRef FName); + } // namespace esimd } // namespace llvm diff --git a/llvm/lib/SYCLLowerIR/ESIMD/ESIMDUtils.cpp b/llvm/lib/SYCLLowerIR/ESIMD/ESIMDUtils.cpp index fe6c1a5509f3e..b7f06fae84918 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/ESIMDUtils.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/ESIMDUtils.cpp @@ -129,6 +129,16 @@ void UpdateUint64MetaDataToMaxValue::operator()(Function *F) const { Node->replaceOperandWith(Key, getMetadata(New)); } } +StringRef stripMangling(StringRef FName) { + + // See if the Name represents an ESIMD intrinsic and demangle only if it + // does. + if (!FName.consume_front(ESIMD_INTRIN_PREF0)) + return ""; + // now skip the digits + FName = FName.drop_while([](char C) { return std::isdigit(C); }); + return FName.starts_with("__esimd") ? FName : ""; +} } // namespace esimd } // namespace llvm diff --git a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp index b0354ff6fdb64..09101f0df7207 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp @@ -130,10 +130,6 @@ enum class lsc_subopcode : uint8_t { read_state_info = 0x1e, fence = 0x1f, }; -// The regexp for ESIMD intrinsics: -// /^_Z(\d+)__esimd_\w+/ -static constexpr char ESIMD_INTRIN_PREF0[] = "_Z"; -static constexpr char ESIMD_INTRIN_PREF1[] = "__esimd_"; static constexpr char ESIMD_INSERTED_VSTORE_FUNC_NAME[] = "_Z14__esimd_vstorev"; static constexpr char SPIRV_INTRIN_PREF[] = "__spirv_BuiltIn"; struct ESIMDIntrinDesc { @@ -2178,12 +2174,11 @@ size_t SYCLLowerESIMDPass::runOnFunction(Function &F, } StringRef Name = Callee->getName(); - // See if the Name represents an ESIMD intrinsic and demangle only if it - // does. - if (!Name.consume_front(ESIMD_INTRIN_PREF0) && !isDevicelibFunction(Name)) + if (!isDevicelibFunction(Name)) + Name = stripMangling(Name); + + if (Name.empty()) continue; - // now skip the digits - Name = Name.drop_while([](char C) { return std::isdigit(C); }); // process ESIMD builtins that go through special handling instead of // the translation procedure diff --git a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp index 013250c3d5bec..3a189304a0272 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMDKernelAttrs.cpp @@ -8,10 +8,11 @@ // Finds and adds sycl_explicit_simd attributes to wrapper functions that wrap // ESIMD kernel functions +#include "llvm/IR/InstIterator.h" +#include "llvm/IR/Module.h" #include "llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h" #include "llvm/SYCLLowerIR/ESIMD/LowerESIMD.h" #include "llvm/SYCLLowerIR/SYCLUtils.h" -#include "llvm/IR/Module.h" #define DEBUG_TYPE "LowerESIMDKernelAttrs" @@ -34,7 +35,20 @@ PreservedAnalyses SYCLFixupESIMDKernelWrapperMDPass::run(Module &M, ModuleAnalysisManager &MAM) { bool Modified = false; for (Function &F : M) { - if (llvm::esimd::isESIMD(F)) { + bool ShouldConsiderESIMD = llvm::esimd::isESIMD(F); + if (!ShouldConsiderESIMD) { + for (Instruction &I : instructions(F)) { + auto *CI = dyn_cast_or_null(&I); + if (!CI) + continue; + auto *CalledF = CI->getCalledFunction(); + if (CalledF && !esimd::stripMangling(CalledF->getName()).empty()) { + ShouldConsiderESIMD = true; + break; + } + } + } + if (ShouldConsiderESIMD) { // TODO: Keep track of traversed functions to avoid repeating traversals // over same function. sycl::utils::traverseCallgraphUp( diff --git a/llvm/test/SYCLLowerIR/ESIMD/prop_metadata.ll b/llvm/test/SYCLLowerIR/ESIMD/prop_metadata.ll new file mode 100644 index 0000000000000..014ee4163155c --- /dev/null +++ b/llvm/test/SYCLLowerIR/ESIMD/prop_metadata.ll @@ -0,0 +1,17 @@ +; This test verifies that we propagate the ESIMD attribute to a function that +; doesn't call any ESIMD-attribute functions but calls an ESIMD intrinsic + +; RUN: opt -passes=lower-esimd-kernel-attrs -S < %s | FileCheck %s + +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" + +; CHECK: define dso_local spir_func void @FUNC() !sycl_explicit_simd +define dso_local spir_func void @FUNC() { + %a_1 = alloca <16 x float> + %1 = load <16 x float>, ptr %a_1 + %ret_val = call spir_func <8 x float> @_Z16__esimd_rdregionIfLi16ELi8ELi0ELi8ELi1ELi0EEN2cm3gen13__vector_typeIT_XT1_EE4typeENS2_IS3_XT0_EE4typeEt(<16 x float> %1, i16 zeroext 0) + ret void +} + +declare dso_local spir_func <8 x float> @_Z16__esimd_rdregionIfLi16ELi8ELi0ELi8ELi1ELi0EEN2cm3gen13__vector_typeIT_XT1_EE4typeENS2_IS3_XT0_EE4typeEt(<16 x float> %0, i16 zeroext %1)