Skip to content

[SYCL][LowerWGScope][SpecConstantsPass] fix new call's calling conv for spir/spirv #15718

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
Oct 21, 2024
Merged

[SYCL][LowerWGScope][SpecConstantsPass] fix new call's calling conv for spir/spirv #15718

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
033a1f6
[SYCL][LowerWGScope] Set new call's calling conv to spirv_func for sp…
wenju-he Oct 16, 2024
c342e0a
isSPIR -> isSPIROrSPIRV
wenju-he Oct 17, 2024
14509a1
fix function and call in SpecConstantsPass
wenju-he Oct 17, 2024
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: 4 additions & 0 deletions llvm/lib/SYCLLowerIR/LowerWGScope.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1001,6 +1001,8 @@ Instruction *spirv::genWGBarrier(Instruction &Before, const Triple &TT) {
FunctionCallee FC =
M.getOrInsertFunction(Name, Attr, RetTy, ScopeTy, ScopeTy, SemanticsTy);
assert(FC.getCallee() && "spirv intrinsic creation failed");
if (TT.isSPIR())
Copy link
Contributor

@sarnex sarnex Oct 16, 2024

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Suggested change
if (TT.isSPIR())
if (TT.isSPIROrSPIRV())

same for other spot too i think

Copy link
Contributor Author

@wenju-he wenju-he Oct 17, 2024

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

done
Also fixed SpecConstantsPass

cast<Function>(FC.getCallee())->setCallingConv(CallingConv::SPIR_FUNC);

IRBuilder<> Bld(Ctx);
Bld.SetInsertPoint(&Before);
Expand All @@ -1011,5 +1013,7 @@ Instruction *spirv::genWGBarrier(Instruction &Before, const Triple &TT) {
asUInt(spirv::MemorySemantics::WorkgroupMemory));
auto BarrierCall = Bld.CreateCall(FC, {ArgExec, ArgMem, ArgSema});
BarrierCall->addFnAttr(llvm::Attribute::Convergent);
if (TT.isSPIR())
BarrierCall->setCallingConv(CallingConv::SPIR_FUNC);
return BarrierCall;
}
27 changes: 27 additions & 0 deletions llvm/test/SYCLLowerIR/LowerWGScope/barrier-calling-conv.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,27 @@
; RUN: opt -passes=LowerWGScope -S %s -o - | FileCheck %s

; Check newly created barrier call has spir_func calling convention.

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-G1"
target triple = "spir64-unknown-unknown"

%"class.sycl::_V1::group" = type { %"class.sycl::_V1::range", %"class.sycl::_V1::range", %"class.sycl::_V1::range", %"class.sycl::_V1::id" }
%"class.sycl::_V1::range" = type { %"class.sycl::_V1::detail::array" }
%"class.sycl::_V1::detail::array" = type { [1 x i64] }
%"class.sycl::_V1::id" = type { %"class.sycl::_V1::detail::array" }

define internal spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_ENKUlNS0_5groupILi1EEEE_clES5_(ptr addrspace(4) noundef align 1 dereferenceable_or_null(1) %this, ptr noundef byval(%"class.sycl::_V1::group") align 8 %group) !work_group_scope !0 {
entry:
; CHECK: call spir_func void @_Z22__spirv_ControlBarrierjjj(

%this.addr = alloca ptr addrspace(4), align 8
%this.addr.ascast = addrspacecast ptr %this.addr to ptr addrspace(4)
store ptr addrspace(4) %this, ptr addrspace(4) %this.addr.ascast, align 8
%group.ascast = addrspacecast ptr %group to ptr addrspace(4)
%this1 = load ptr addrspace(4), ptr addrspace(4) %this.addr.ascast, align 8
ret void
}

; CHECK: declare spir_func void @_Z22__spirv_ControlBarrierjjj(

!0 = !{}
Loading