Skip to content

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

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 3 commits into from
Apr 13, 2021
Merged

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

Show file tree
Hide file tree
Changes from 2 commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
f602443
[ESIMD] Skip rewriting functions used through function pointers
DenisBakhvalov Apr 9, 2021
86b638d
Fixed a line width
DenisBakhvalov Apr 9, 2021
343e310
Added a comment about skip rewriting
DenisBakhvalov Apr 12, 2021
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 3 additions & 1 deletion llvm/lib/SYCLLowerIR/LowerESIMDVecArg.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -255,7 +255,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) {
Expand Down
59 changes: 59 additions & 0 deletions llvm/test/SYCLLowerIR/esimd_lower_vec_arg_fp.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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 = !{}