[SYCL][ESIMD] Fix propagation of ESIMD attribute for inlined functions

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

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

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

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<CallInst>(&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(

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)