Skip to content

Commit 0ae2950

Browse files
authored
[SYCL] Support passing bool as kernel param for SPIR-V target (#17427)
When targeting SPIR-V, the compiler was incorrectly lowering bool to i1. This caused the translator to produce an OpTypeBool which is illegal for an entry point in SPIR-V. Instead, the patch coerce the type to follow the host size. Fix #11531
1 parent 639c897 commit 0ae2950

File tree

3 files changed

+69
-2
lines changed

3 files changed

+69
-2
lines changed

clang/lib/CodeGen/Targets/SPIR.cpp

Lines changed: 13 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -38,9 +38,20 @@ class CommonSPIRABIInfo : public DefaultABIInfo {
3838
ABIArgInfo CommonSPIRABIInfo::classifyKernelArgumentType(QualType Ty) const {
3939
Ty = useFirstFieldIfTransparentUnion(Ty);
4040

41-
if (getContext().getLangOpts().SYCLIsDevice && isAggregateTypeForABI(Ty)) {
41+
if (getContext().getLangOpts().SYCLIsDevice) {
42+
if (const BuiltinType *BT = Ty->getAs<BuiltinType>()) {
43+
switch (BT->getKind()) {
44+
case BuiltinType::Bool:
45+
// Bool / i1 isn't a legal kernel argument in SPIR-V.
46+
// Coerce the type to follow the host representation of bool.
47+
return ABIArgInfo::getDirect(CGT.ConvertTypeForMem(Ty));
48+
default:
49+
break;
50+
}
51+
}
4252
// Pass all aggregate types allowed by Sema by value.
43-
return getNaturalAlignIndirect(Ty);
53+
if (isAggregateTypeForABI(Ty))
54+
return getNaturalAlignIndirect(Ty);
4455
}
4556

4657
return DefaultABIInfo::classifyArgumentType(Ty);
Lines changed: 26 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,26 @@
1+
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
2+
3+
#include "sycl.hpp"
4+
5+
void take_bool(bool) {}
6+
7+
int main() {
8+
bool test = false;
9+
sycl::queue q;
10+
11+
// CHECK: @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E11test_kernel(i8 {{.*}} [[ARG:%[A-Za-z_0-9]*]]
12+
// CHECK: %__SYCLKernel = alloca
13+
// CHECK: %test = getelementptr inbounds nuw %class.anon, ptr addrspace(4) %__SYCLKernel.ascast
14+
// CHECK: store i8 %{{.*}}, ptr addrspace(4) %test
15+
// CHECK: call spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_ENKUlvE_clEv
16+
//
17+
// CHECK: define {{.*}} @_Z9take_boolb(i1
18+
q.submit([&](sycl::handler &h) {
19+
h.single_task<class test_kernel>([=]() {
20+
(void)test;
21+
take_bool(test);
22+
});
23+
});
24+
25+
return 0;
26+
}

sycl/test-e2e/Basic/bool-test.cpp

Lines changed: 30 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,30 @@
1+
// RUN: %{build} -o %t.out
2+
// RUN: %{run} %t.out
3+
4+
// Check booleans are promoted correctly
5+
6+
#include <sycl/detail/core.hpp>
7+
#include <sycl/usm.hpp>
8+
9+
void run_test(sycl::queue q, bool test, int *res) {
10+
q.submit([&](sycl::handler &cgh) {
11+
cgh.single_task([=]() {
12+
if (test)
13+
*res = 42;
14+
else
15+
*res = -42;
16+
});
17+
}).wait();
18+
}
19+
20+
int main() {
21+
sycl::queue q;
22+
int *p = sycl::malloc_shared<int>(1, q);
23+
*p = 0;
24+
run_test(q, true, p);
25+
assert(*p == 42);
26+
*p = 0;
27+
run_test(q, false, p);
28+
assert(*p == -42);
29+
sycl::free(p, q);
30+
}

0 commit comments

Comments
 (0)