[SYCL][ESIMD] Fix propagation of ESIMD attribute for inlined functions (#16193)

The internal compiler uses aggressive inlining, and we have a case where
a wrapper function that calls an ESIMD function has had the ESIMD
functions body inlined into it. The parent function does not have the
ESIMD attribute, so this causes the pass that propagates the ESIMD
attribute to fail, which ends up causing `sycl-post-link` to do the
wrong thing. It works fine when the inlining does not happen because
this pass propagates the attribute. Extend the ESIMD attribute
propagation pass to also consider functions that call ESIMD intrinsics
as ESIMD functions.

---------

Signed-off-by: Sarnie, Nick <[email protected]>

Author: sarnex

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 mangling 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)