From f602443b4e308f2c38ab9883dc7d374c6468906d Mon Sep 17 00:00:00 2001 From: Denis Bakhvalov Date: Fri, 9 Apr 2021 14:14:54 -0700 Subject: [PATCH 1/3] [ESIMD] Skip rewriting functions used through function pointers --- llvm/lib/SYCLLowerIR/LowerESIMDVecArg.cpp | 4 +- .../SYCLLowerIR/esimd_lower_vec_arg_fp.ll | 58 +++++++++++++++++++ 2 files changed, 61 insertions(+), 1 deletion(-) create mode 100644 llvm/test/SYCLLowerIR/esimd_lower_vec_arg_fp.ll diff --git a/llvm/lib/SYCLLowerIR/LowerESIMDVecArg.cpp b/llvm/lib/SYCLLowerIR/LowerESIMDVecArg.cpp index 8d99cb219c0f3..c6f1460884c8b 100644 --- a/llvm/lib/SYCLLowerIR/LowerESIMDVecArg.cpp +++ b/llvm/lib/SYCLLowerIR/LowerESIMDVecArg.cpp @@ -255,7 +255,9 @@ PreservedAnalyses ESIMDLowerVecArgPass::run(Module &M, SmallVector 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..938b25bd4d9af --- /dev/null +++ b/llvm/test/SYCLLowerIR/esimd_lower_vec_arg_fp.ll @@ -0,0 +1,58 @@ +; 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 = !{} From 86b638d26a88384fdf4196afaab0b3f862b611d4 Mon Sep 17 00:00:00 2001 From: Denis Bakhvalov Date: Fri, 9 Apr 2021 14:17:30 -0700 Subject: [PATCH 2/3] Fixed a line width --- llvm/test/SYCLLowerIR/esimd_lower_vec_arg_fp.ll | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/llvm/test/SYCLLowerIR/esimd_lower_vec_arg_fp.ll b/llvm/test/SYCLLowerIR/esimd_lower_vec_arg_fp.ll index 938b25bd4d9af..2ba4973107cc0 100644 --- a/llvm/test/SYCLLowerIR/esimd_lower_vec_arg_fp.ll +++ b/llvm/test/SYCLLowerIR/esimd_lower_vec_arg_fp.ll @@ -1,7 +1,8 @@ ; 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. +; 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> } From 343e3100f0ae696b597dc74f164704a32019d945 Mon Sep 17 00:00:00 2001 From: Denis Bakhvalov Date: Mon, 12 Apr 2021 13:10:11 -0700 Subject: [PATCH 3/3] Added a comment about skip rewriting --- llvm/lib/SYCLLowerIR/LowerESIMDVecArg.cpp | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/llvm/lib/SYCLLowerIR/LowerESIMDVecArg.cpp b/llvm/lib/SYCLLowerIR/LowerESIMDVecArg.cpp index c6f1460884c8b..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,7 @@ PreservedAnalyses ESIMDLowerVecArgPass::run(Module &M, SmallVector functions; for (auto &F : M) { - // Skip functions that are used through function pointers + // Skip functions that are used through function pointers. if (!F.hasAddressTaken()) functions.push_back(&F); }