diff --git a/llvm/lib/SYCLLowerIR/LowerESIMDVecArg.cpp b/llvm/lib/SYCLLowerIR/LowerESIMDVecArg.cpp
index 8d99cb219c0f3..9b30c72d3436a 100644
--- a/llvm/lib/SYCLLowerIR/LowerESIMDVecArg.cpp
+++ b/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) {
diff --git a/llvm/test/SYCLLowerIR/esimd_lower_vec_arg_fp.ll b/llvm/test/SYCLLowerIR/esimd_lower_vec_arg_fp.ll
new file mode 100644
index 0000000000000..2ba4973107cc0
--- /dev/null
+++ b/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 = !{}