Skip to content

Commit d4478a4

Browse files
elizabethandrewselizabethandrews
authored
[SYCL] Do not decompose arrays with pointers (#7015)
Arrays with pointers are no longer decomposed into array elements when generating openCL kernel arguments. Signed-off-by: Elizabeth Andrews <[email protected]>
1 parent 3c7afff commit d4478a4

File tree

8 files changed

+451
-447
lines changed

8 files changed

+451
-447
lines changed

‎clang/lib/Sema/SemaSYCL.cpp‎

Lines changed: 212 additions & 107 deletions
Large diffs are not rendered by default.

‎clang/test/CodeGenSYCL/inheritance.cpp‎

Lines changed: 4 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -5,6 +5,7 @@
55
class second_base {
66
public:
77
int *e;
8+
int *arr[2];
89
second_base(int *E) : e(E) {}
910
};
1011

@@ -43,9 +44,9 @@ int main() {
4344
// CHECK: %struct.base = type { i32, %class.InnerField }
4445
// CHECK: %class.InnerField = type { %class.InnerFieldBase, i32 }
4546
// CHECK: %class.InnerFieldBase = type { i32 }
46-
// CHECK: %class.__generated_second_base = type { ptr addrspace(1) }
47+
// CHECK: %class.__generated_second_base = type { ptr addrspace(1), [2 x ptr addrspace(1)] }
4748
// CHECK: %struct.derived = type <{ %struct.base, [4 x i8], %class.second_base, i32, [4 x i8] }>
48-
// CHECK: %class.second_base = type { ptr addrspace(4) }
49+
// CHECK: %class.second_base = type { ptr addrspace(4), [2 x ptr addrspace(4)] }
4950

5051
// Check kernel paramters
5152
// CHECK: define {{.*}}spir_kernel void @{{.*}}derived
@@ -69,7 +70,7 @@ int main() {
6970
// First, derived-to-base cast with offset:
7071
// CHECK: %[[OFFSET_CALC:.*]] = getelementptr inbounds i8, ptr addrspace(4) %[[LOCAL_OBJECT]], i64 16
7172
// Initialize 'second_base'
72-
// CHECK: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 8 %[[OFFSET_CALC]], ptr addrspace(4) align 8 %[[ARG_BASE1]], i64 8, i1 false)
73+
// CHECK: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 8 %[[OFFSET_CALC]], ptr addrspace(4) align 8 %[[ARG_BASE1]], i64 24, i1 false)
7374

7475
// Initialize field 'a'
7576
// CHECK: %[[GEP_A:[a-zA-Z0-9]+]] = getelementptr inbounds %struct.derived, ptr addrspace(4) %[[LOCAL_OBJECT]], i32 0, i32 3

‎clang/test/CodeGenSYCL/no_opaque_pointers-in-structs.cpp‎

Lines changed: 3 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -33,16 +33,9 @@ int main() {
3333
return 0;
3434
}
3535

36-
// CHECK: %[[WRAPPER_F1:[a-zA-Z0-9_.]+]] = type { i32 addrspace(1)* }
37-
// CHECK: %[[WRAPPER_F2:[a-zA-Z0-9_.]+]] = type { float addrspace(1)* }
38-
// CHECK: %[[GENERATED_A:[a-zA-Z0-9_.]+]] = type { float addrspace(1)* }
39-
// CHECK: %[[WRAPPER_F4_1:[a-zA-Z0-9_.]+]] = type { i32 addrspace(1)* }
40-
// CHECK: %[[WRAPPER_F4_2:[a-zA-Z0-9_.]+]] = type { i32 addrspace(1)* }
36+
// CHECK: %[[GENERATED_B:[a-zA-Z0-9_.]+]] = type { i32 addrspace(1)*, float addrspace(1)*, %[[GENERATED_A:[a-zA-Z0-9_.]+]], [2 x i32 addrspace(1)*] }
37+
// CHECK: %[[GENERATED_A]] = type { float addrspace(1)* }
4138
// CHECK: %[[WRAPPER_LAMBDA_PTR:[a-zA-Z0-9_.]+]] = type { float addrspace(1)* }
4239
// CHECK: define {{.*}}spir_kernel void @{{.*}}structs
43-
// CHECK-SAME: %[[WRAPPER_F1]]* noundef byval(%[[WRAPPER_F1]]) align 8 %_arg_F1,
44-
// CHECK-SAME: %[[WRAPPER_F2]]* noundef byval(%[[WRAPPER_F2]]) align 8 %_arg_F2,
45-
// CHECK-SAME: %[[GENERATED_A]]* noundef byval(%[[GENERATED_A]]) align 8 %_arg_F3,
46-
// CHECK-SAME: %[[WRAPPER_F4_1]]* noundef byval(%[[WRAPPER_F4_1]]) align 8 %_arg_F4
47-
// CHECK-SAME: %[[WRAPPER_F4_2]]* noundef byval(%[[WRAPPER_F4_2]]) align 8 %_arg_F41
40+
// CHECK-SAME: %[[GENERATED_B]]* noundef byval(%[[GENERATED_B]]) align 8 %_arg_Obj
4841
// CHECK: define {{.*}}spir_kernel void @{{.*}}lambdas{{.*}}(%[[WRAPPER_LAMBDA_PTR]]* noundef byval(%[[WRAPPER_LAMBDA_PTR]]) align 8 %_arg_Lambda)

‎clang/test/CodeGenSYCL/pointers-in-structs.cpp‎

Lines changed: 4 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -33,16 +33,10 @@ int main() {
3333
return 0;
3434
}
3535

36-
// CHECK: %[[WRAPPER_F1:[a-zA-Z0-9_.]+]] = type { ptr addrspace(1) }
37-
// CHECK: %[[WRAPPER_F2:[a-zA-Z0-9_.]+]] = type { ptr addrspace(1) }
38-
// CHECK: %[[GENERATED_A:[a-zA-Z0-9_.]+]] = type { ptr addrspace(1) }
39-
// CHECK: %[[WRAPPER_F4_1:[a-zA-Z0-9_.]+]] = type { ptr addrspace(1) }
40-
// CHECK: %[[WRAPPER_F4_2:[a-zA-Z0-9_.]+]] = type { ptr addrspace(1) }
36+
// CHECK: %[[GENERATED_B:[a-zA-Z0-9_.]+]] = type { ptr addrspace(1), ptr addrspace(1), %[[GENERATED_A:[a-zA-Z0-9_.]+]], [2 x ptr addrspace(1)] }
37+
// CHECK: [[GENERATED_A]] = type { ptr addrspace(1) }
4138
// CHECK: %[[WRAPPER_LAMBDA_PTR:[a-zA-Z0-9_.]+]] = type { ptr addrspace(1) }
39+
4240
// CHECK: define {{.*}}spir_kernel void @{{.*}}structs
43-
// CHECK-SAME: ptr noundef byval(%[[WRAPPER_F1]]) align 8 %_arg_F1,
44-
// CHECK-SAME: ptr noundef byval(%[[WRAPPER_F2]]) align 8 %_arg_F2,
45-
// CHECK-SAME: ptr noundef byval(%[[GENERATED_A]]) align 8 %_arg_F3,
46-
// CHECK-SAME: ptr noundef byval(%[[WRAPPER_F4_1]]) align 8 %_arg_F4
47-
// CHECK-SAME: ptr noundef byval(%[[WRAPPER_F4_2]]) align 8 %_arg_F41
41+
// CHECK-SAME: ptr noundef byval(%[[GENERATED_B]]) align 8 %_arg_Obj
4842
// CHECK: define {{.*}}spir_kernel void @{{.*}}lambdas{{.*}}(ptr noundef byval(%[[WRAPPER_LAMBDA_PTR]]) align 8 %_arg_Lambda)

‎clang/test/CodeGenSYCL/pointers-int-header.cpp‎

Lines changed: 18 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -7,7 +7,15 @@
77

88
#include "Inputs/sycl.hpp"
99

10-
struct struct_with_pointer {
10+
struct decomposed_struct_with_pointer {
11+
int data_in_struct;
12+
int *ptr_in_struct;
13+
int *ptr_array_in_struct1[2];
14+
int *ptr_array_in_struct2[2][3];
15+
sycl::accessor<char, 1, sycl::access::mode::read> acc;
16+
};
17+
18+
struct non_decomposed_struct_with_pointer {
1119
int data_in_struct;
1220
int *ptr_in_struct;
1321
int *ptr_array_in_struct1[2];
@@ -16,24 +24,22 @@ struct struct_with_pointer {
1624

1725
int main() {
1826
int *ptr;
19-
struct_with_pointer obj;
20-
obj.data_in_struct = 10;
27+
decomposed_struct_with_pointer obj1;
28+
non_decomposed_struct_with_pointer obj2;
29+
obj1.data_in_struct = 10;
30+
obj2.data_in_struct = 10;
2131

2232
sycl::kernel_single_task<class test>([=]() {
2333
*ptr = 50;
24-
int local = obj.data_in_struct;
34+
int local = obj1.data_in_struct + obj2.data_in_struct;
2535
});
2636
}
2737

2838
// Integration header entries for pointer, scalar and wrapped pointer.
2939
// CHECK:{ kernel_param_kind_t::kind_pointer, 8, 0 },
3040
// CHECK:{ kernel_param_kind_t::kind_std_layout, 4, 8 },
3141
// CHECK:{ kernel_param_kind_t::kind_std_layout, 8, 16 },
32-
// CHECK:{ kernel_param_kind_t::kind_std_layout, 8, 24 },
33-
// CHECK:{ kernel_param_kind_t::kind_std_layout, 8, 32 },
34-
// CHECK:{ kernel_param_kind_t::kind_std_layout, 8, 40 },
35-
// CHECK:{ kernel_param_kind_t::kind_std_layout, 8, 48 },
36-
// CHECK:{ kernel_param_kind_t::kind_std_layout, 8, 56 },
37-
// CHECK:{ kernel_param_kind_t::kind_std_layout, 8, 64 },
38-
// CHECK:{ kernel_param_kind_t::kind_std_layout, 8, 72 },
39-
// CHECK:{ kernel_param_kind_t::kind_std_layout, 8, 80 },
42+
// CHECK:{ kernel_param_kind_t::kind_std_layout, 16, 24 },
43+
// CHECK:{ kernel_param_kind_t::kind_std_layout, 48, 40 },
44+
// CHECK:{ kernel_param_kind_t::kind_accessor, 4062, 88 },
45+
// CHECK:{ kernel_param_kind_t::kind_std_layout, 80, 104 },

0 commit comments

Comments
 (0)