Skip to content

Commit 449308d

Browse files
authored
[SYCL] Add a test for wrapped pointer kernel args (#2298)
Before structs decomposition USM pointers that are wrapped by structs were passsed with generic address space. It caused performance issues. After structs decomposition all pointers are passed with global address space, this test is added to save this behaviour.
1 parent 43862a3 commit 449308d

File tree

1 file changed

+36
-0
lines changed

1 file changed

+36
-0
lines changed
Lines changed: 36 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,36 @@
1+
// RUN: %clang_cc1 -fsycl -fsycl-is-device -I %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
2+
3+
// This test checks that compiler generates correct address spaces for pointer
4+
// kernel arguments that are wrapped by struct.
5+
6+
#include <sycl.hpp>
7+
8+
struct A {
9+
float *F;
10+
};
11+
12+
struct B {
13+
int *F1;
14+
float *F2;
15+
A F3;
16+
};
17+
18+
int main() {
19+
B Obj;
20+
cl::sycl::kernel_single_task<class structs>(
21+
[=]() {
22+
(void)Obj;
23+
});
24+
float A = 1;
25+
float *Ptr = &A;
26+
auto Lambda = [=]() {
27+
*Ptr += 1;
28+
};
29+
cl::sycl::kernel_single_task<class lambdas>([=]() {
30+
Lambda();
31+
});
32+
return 0;
33+
}
34+
35+
// CHECK: define spir_kernel void @{{.*}}structs{{.*}}(i32 addrspace(1)* %_arg_F1, float addrspace(1)* %_arg_F2, float addrspace(1)* %_arg_F)
36+
// CHECK: define spir_kernel void @{{.*}}lambdas{{.*}}(float addrspace(1)* %_arg_)

0 commit comments

Comments
 (0)