[ESIMD] Skip rewriting functions used through function pointers (#3527)

This fixes a crash when running one of the ESIMD tests that use function pointers with `-O0`. The problem here is that the `ESIMDLowerVecArg` pass expects users of the functions to be only call instructions, but it's not the case with function pointers since the address of a function can be stored in some variable.

Author: DenisBakhvalov

llvm/lib/SYCLLowerIR/LowerESIMDVecArg.cpp

@@ -31,6 +31,9 @@
 // % 1 = bitcast<16 x i32> * % 0 to %
 // "class._ZTSN2cm3gen4simdIiLi16EEE.cm::gen::simd" *
 //
+// It is OK not to rewrite a function (for example, when its address is taken)
+// since it does not affect correctness. But that may lead to vector backend
+// not being able to hold the value in GRF and generate memory references.
 //
 // Change in global variables:
 //
@@ -255,7 +258,9 @@ PreservedAnalyses ESIMDLowerVecArgPass::run(Module &M,
 
   SmallVector<Function *, 10> functions;
   for (auto &F : M) {
-    functions.push_back(&F);
+    // Skip functions that are used through function pointers.
+    if (!F.hasAddressTaken())
+      functions.push_back(&F);
   }
 
   for (auto F : functions) {

llvm/test/SYCLLowerIR/esimd_lower_vec_arg_fp.ll

@@ -0,0 +1,59 @@
+; NOTE: Assertions have been autogenerated by utils/update_test_checks.py
+; RUN: opt < %s -ESIMDLowerVecArg -S | FileCheck %s
+
+; This test checks that there is no crash in ESIMDLowerVecArg pass when
+; rewriting funcitons that are used through a function pointer.
+
+%"cl::sycl::INTEL::gpu::simd" = type { <64 x i32> }
+
+define dso_local spir_func void @func(%"cl::sycl::INTEL::gpu::simd"* %arg) {
+; CHECK-LABEL: @func(
+; CHECK-NEXT:  entry:
+; CHECK-NEXT:    ret void
+;
+entry:
+  ret void
+}
+
+define dso_local spir_func void @init_ptr(void (%"cl::sycl::INTEL::gpu::simd"*)** %foo) !sycl_explicit_simd !1 {
+; CHECK-LABEL: @init_ptr(
+; CHECK-NEXT:  entry:
+; CHECK-NEXT:    store void (%"cl::sycl::INTEL::gpu::simd"*)* @func, void (%"cl::sycl::INTEL::gpu::simd"*)** [[FOO:%.*]], align 8
+; CHECK-NEXT:    ret void
+;
+entry:
+  store void (%"cl::sycl::INTEL::gpu::simd"*)* @func, void (%"cl::sycl::INTEL::gpu::simd"*)** %foo
+  ret void
+}
+
+define dso_local spir_func void @use_ptr(void (%"cl::sycl::INTEL::gpu::simd"*)* %foo) !sycl_explicit_simd !1 {
+; CHECK-LABEL: @use_ptr(
+; CHECK-NEXT:  entry:
+; CHECK-NEXT:    [[AGG_TMP:%.*]] = alloca %"cl::sycl::INTEL::gpu::simd", align 256
+; CHECK-NEXT:    call spir_func void [[FOO:%.*]](%"cl::sycl::INTEL::gpu::simd"* [[AGG_TMP]])
+; CHECK-NEXT:    ret void
+;
+entry:
+  %agg.tmp = alloca %"cl::sycl::INTEL::gpu::simd"
+  call spir_func void %foo(%"cl::sycl::INTEL::gpu::simd"* %agg.tmp)
+  ret void
+}
+
+define dso_local spir_func void @esimd_kernel() !sycl_explicit_simd !1 {
+; CHECK-LABEL: @esimd_kernel(
+; CHECK-NEXT:  entry:
+; CHECK-NEXT:    [[FP:%.*]] = alloca void (%"cl::sycl::INTEL::gpu::simd"*)*, align 8
+; CHECK-NEXT:    call spir_func void @init_ptr(void (%"cl::sycl::INTEL::gpu::simd"*)** [[FP]])
+; CHECK-NEXT:    [[TMP0:%.*]] = load void (%"cl::sycl::INTEL::gpu::simd"*)*, void (%"cl::sycl::INTEL::gpu::simd"*)** [[FP]], align 8
+; CHECK-NEXT:    call spir_func void @use_ptr(void (%"cl::sycl::INTEL::gpu::simd"*)* [[TMP0]])
+; CHECK-NEXT:    ret void
+;
+entry:
+  %fp = alloca void (%"cl::sycl::INTEL::gpu::simd"*)*
+  call spir_func void @init_ptr(void (%"cl::sycl::INTEL::gpu::simd"*)** %fp)
+  %0 = load void (%"cl::sycl::INTEL::gpu::simd"*)*, void (%"cl::sycl::INTEL::gpu::simd"*)** %fp
+  call spir_func void @use_ptr(void (%"cl::sycl::INTEL::gpu::simd"*)* %0)
+  ret void
+}
+
+!1 = !{}