Skip to content

Commit 75edca0

Browse files
committed
Apply initial feedback
1 parent 2340c2c commit 75edca0

File tree

2 files changed

+23
-3
lines changed

2 files changed

+23
-3
lines changed

llvm/lib/SYCLLowerIR/LowerWGLocalMemory.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -71,10 +71,10 @@ sycl::getKernelNamesUsingImplicitLocalMem(const Module &M) {
7171
return -1;
7272
};
7373
llvm::for_each(M.functions(), [&](const Function &F) {
74-
if (F.getCallingConv() == CallingConv::SPIR_KERNEL) {
74+
if (F.getCallingConv() == CallingConv::SPIR_KERNEL &&
75+
F.hasFnAttribute(WORK_GROUP_STATIC_ATTR)) {
7576
int ArgPos = GetArgumentPos(F);
76-
if (ArgPos >= 0 || F.hasFnAttribute(WORK_GROUP_STATIC_ATTR))
77-
SPIRKernelNames.emplace_back(F.getName(), ArgPos);
77+
SPIRKernelNames.emplace_back(F.getName(), ArgPos);
7878
}
7979
});
8080
}

llvm/test/SYCLLowerIR/work_group_static.ll

Lines changed: 20 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -22,9 +22,29 @@ entry:
2222
ret void
2323
}
2424

25+
; Function Attrs: convergent norecurse
26+
; CHECK: @__sycl_kernel_B{{.*}} #[[ATTRS:[0-9]+]]
27+
define weak_odr dso_local spir_kernel void @__sycl_kernel_B(ptr addrspace(1) %0) local_unnamed_addr #1 !kernel_arg_addr_space !5 {
28+
entry:
29+
%1 = tail call spir_func ptr addrspace(3) @__sycl_dynamicLocalMemoryPlaceholder(i64 128) #1
30+
ret void
31+
}
32+
33+
; Function Attrs: convergent norecurse
34+
; CHECK: @__sycl_kernel_C{{.*}} #[[ATTRS]]
35+
define weak_odr dso_local spir_kernel void @__sycl_kernel_C(ptr addrspace(1) %0) local_unnamed_addr #1 !kernel_arg_addr_space !5 {
36+
entry:
37+
%1 = tail call spir_func ptr addrspace(3) @__sycl_allocateLocalMemory(i64 128, i64 4) #1
38+
ret void
39+
}
40+
41+
; Function Attrs: convergent
42+
declare dso_local spir_func ptr addrspace(3) @__sycl_allocateLocalMemory(i64, i64) local_unnamed_addr #1
43+
2544
; Function Attrs: convergent
2645
declare dso_local spir_func ptr addrspace(3) @__sycl_dynamicLocalMemoryPlaceholder(i64) local_unnamed_addr #1
2746

47+
; CHECK: #[[ATTRS]] = {{.*}} "sycl-work-group-static"
2848
attributes #0 = { convergent norecurse "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" "sycl-work-group-static"="1" }
2949
attributes #1 = { convergent norecurse }
3050

0 commit comments

Comments
 (0)