Skip to content

[SYCL] Do not decompose arrays with pointers #7015

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 6 commits into from
Oct 13, 2022
Merged

[SYCL] Do not decompose arrays with pointers #7015

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
43c9dd0
[SYCL] Do not decompose arrays with pointers
elizabethandrews Oct 9, 2022
2dfe874
Apply review comments
elizabethandrews Oct 11, 2022
28a35da
Apply review comments
elizabethandrews Oct 12, 2022
d60780a
Apply review comment
elizabethandrews Oct 12, 2022
81dfe20
Merge remote-tracking branch 'intel_llvm/sycl' into DontDecomposeArray
elizabethandrews Oct 13, 2022
83930d7
Merge remote-tracking branch 'intel_llvm/sycl' into DontDecomposeArray
elizabethandrews Oct 13, 2022
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
319 changes: 212 additions & 107 deletions clang/lib/Sema/SemaSYCL.cpp
View file
Open in desktop

Large diffs are not rendered by default.

7 changes: 4 additions & 3 deletions clang/test/CodeGenSYCL/inheritance.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,6 +5,7 @@
class second_base {
public:
int *e;
int *arr[2];
second_base(int *E) : e(E) {}
};

Expand Down Expand Up @@ -43,9 +44,9 @@ int main() {
// CHECK: %struct.base = type { i32, %class.InnerField }
// CHECK: %class.InnerField = type { %class.InnerFieldBase, i32 }
// CHECK: %class.InnerFieldBase = type { i32 }
// CHECK: %class.__generated_second_base = type { ptr addrspace(1) }
// CHECK: %class.__generated_second_base = type { ptr addrspace(1), [2 x ptr addrspace(1)] }
// CHECK: %struct.derived = type <{ %struct.base, [4 x i8], %class.second_base, i32, [4 x i8] }>
// CHECK: %class.second_base = type { ptr addrspace(4) }
// CHECK: %class.second_base = type { ptr addrspace(4), [2 x ptr addrspace(4)] }

// Check kernel paramters
// CHECK: define {{.*}}spir_kernel void @{{.*}}derived
Expand All @@ -69,7 +70,7 @@ int main() {
// First, derived-to-base cast with offset:
// CHECK: %[[OFFSET_CALC:.*]] = getelementptr inbounds i8, ptr addrspace(4) %[[LOCAL_OBJECT]], i64 16
// Initialize 'second_base'
// 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)
// 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)

// Initialize field 'a'
// CHECK: %[[GEP_A:[a-zA-Z0-9]+]] = getelementptr inbounds %struct.derived, ptr addrspace(4) %[[LOCAL_OBJECT]], i32 0, i32 3
Expand Down
13 changes: 3 additions & 10 deletions clang/test/CodeGenSYCL/no_opaque_pointers-in-structs.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -33,16 +33,9 @@ int main() {
return 0;
}

// CHECK: %[[WRAPPER_F1:[a-zA-Z0-9_.]+]] = type { i32 addrspace(1)* }
// CHECK: %[[WRAPPER_F2:[a-zA-Z0-9_.]+]] = type { float addrspace(1)* }
// CHECK: %[[GENERATED_A:[a-zA-Z0-9_.]+]] = type { float addrspace(1)* }
// CHECK: %[[WRAPPER_F4_1:[a-zA-Z0-9_.]+]] = type { i32 addrspace(1)* }
// CHECK: %[[WRAPPER_F4_2:[a-zA-Z0-9_.]+]] = type { i32 addrspace(1)* }
// 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)*] }
// CHECK: %[[GENERATED_A]] = type { float addrspace(1)* }
// CHECK: %[[WRAPPER_LAMBDA_PTR:[a-zA-Z0-9_.]+]] = type { float addrspace(1)* }
// CHECK: define {{.*}}spir_kernel void @{{.*}}structs
// CHECK-SAME: %[[WRAPPER_F1]]* noundef byval(%[[WRAPPER_F1]]) align 8 %_arg_F1,
// CHECK-SAME: %[[WRAPPER_F2]]* noundef byval(%[[WRAPPER_F2]]) align 8 %_arg_F2,
// CHECK-SAME: %[[GENERATED_A]]* noundef byval(%[[GENERATED_A]]) align 8 %_arg_F3,
// CHECK-SAME: %[[WRAPPER_F4_1]]* noundef byval(%[[WRAPPER_F4_1]]) align 8 %_arg_F4
// CHECK-SAME: %[[WRAPPER_F4_2]]* noundef byval(%[[WRAPPER_F4_2]]) align 8 %_arg_F41
// CHECK-SAME: %[[GENERATED_B]]* noundef byval(%[[GENERATED_B]]) align 8 %_arg_Obj
// CHECK: define {{.*}}spir_kernel void @{{.*}}lambdas{{.*}}(%[[WRAPPER_LAMBDA_PTR]]* noundef byval(%[[WRAPPER_LAMBDA_PTR]]) align 8 %_arg_Lambda)
14 changes: 4 additions & 10 deletions clang/test/CodeGenSYCL/pointers-in-structs.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -33,16 +33,10 @@ int main() {
return 0;
}

// CHECK: %[[WRAPPER_F1:[a-zA-Z0-9_.]+]] = type { ptr addrspace(1) }
// CHECK: %[[WRAPPER_F2:[a-zA-Z0-9_.]+]] = type { ptr addrspace(1) }
// CHECK: %[[GENERATED_A:[a-zA-Z0-9_.]+]] = type { ptr addrspace(1) }
// CHECK: %[[WRAPPER_F4_1:[a-zA-Z0-9_.]+]] = type { ptr addrspace(1) }
// CHECK: %[[WRAPPER_F4_2:[a-zA-Z0-9_.]+]] = type { ptr addrspace(1) }
// 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)] }
// CHECK: [[GENERATED_A]] = type { ptr addrspace(1) }
// CHECK: %[[WRAPPER_LAMBDA_PTR:[a-zA-Z0-9_.]+]] = type { ptr addrspace(1) }

// CHECK: define {{.*}}spir_kernel void @{{.*}}structs
// CHECK-SAME: ptr noundef byval(%[[WRAPPER_F1]]) align 8 %_arg_F1,
// CHECK-SAME: ptr noundef byval(%[[WRAPPER_F2]]) align 8 %_arg_F2,
// CHECK-SAME: ptr noundef byval(%[[GENERATED_A]]) align 8 %_arg_F3,
// CHECK-SAME: ptr noundef byval(%[[WRAPPER_F4_1]]) align 8 %_arg_F4
// CHECK-SAME: ptr noundef byval(%[[WRAPPER_F4_2]]) align 8 %_arg_F41
// CHECK-SAME: ptr noundef byval(%[[GENERATED_B]]) align 8 %_arg_Obj
// CHECK: define {{.*}}spir_kernel void @{{.*}}lambdas{{.*}}(ptr noundef byval(%[[WRAPPER_LAMBDA_PTR]]) align 8 %_arg_Lambda)
30 changes: 18 additions & 12 deletions clang/test/CodeGenSYCL/pointers-int-header.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -7,7 +7,15 @@

#include "Inputs/sycl.hpp"

struct struct_with_pointer {
struct decomposed_struct_with_pointer {
int data_in_struct;
int *ptr_in_struct;
int *ptr_array_in_struct1[2];
int *ptr_array_in_struct2[2][3];
sycl::accessor<char, 1, sycl::access::mode::read> acc;
};

struct non_decomposed_struct_with_pointer {
int data_in_struct;
int *ptr_in_struct;
int *ptr_array_in_struct1[2];
Expand All @@ -16,24 +24,22 @@ struct struct_with_pointer {

int main() {
int *ptr;
struct_with_pointer obj;
obj.data_in_struct = 10;
decomposed_struct_with_pointer obj1;
non_decomposed_struct_with_pointer obj2;
obj1.data_in_struct = 10;
obj2.data_in_struct = 10;

sycl::kernel_single_task<class test>([=]() {
*ptr = 50;
int local = obj.data_in_struct;
int local = obj1.data_in_struct + obj2.data_in_struct;
});
}

// Integration header entries for pointer, scalar and wrapped pointer.
// CHECK:{ kernel_param_kind_t::kind_pointer, 8, 0 },
// CHECK:{ kernel_param_kind_t::kind_std_layout, 4, 8 },
// CHECK:{ kernel_param_kind_t::kind_std_layout, 8, 16 },
// CHECK:{ kernel_param_kind_t::kind_std_layout, 8, 24 },
// CHECK:{ kernel_param_kind_t::kind_std_layout, 8, 32 },
// CHECK:{ kernel_param_kind_t::kind_std_layout, 8, 40 },
// CHECK:{ kernel_param_kind_t::kind_std_layout, 8, 48 },
// CHECK:{ kernel_param_kind_t::kind_std_layout, 8, 56 },
// CHECK:{ kernel_param_kind_t::kind_std_layout, 8, 64 },
// CHECK:{ kernel_param_kind_t::kind_std_layout, 8, 72 },
// CHECK:{ kernel_param_kind_t::kind_std_layout, 8, 80 },
// CHECK:{ kernel_param_kind_t::kind_std_layout, 16, 24 },
// CHECK:{ kernel_param_kind_t::kind_std_layout, 48, 40 },
// CHECK:{ kernel_param_kind_t::kind_accessor, 4062, 88 },
// CHECK:{ kernel_param_kind_t::kind_std_layout, 80, 104 },
Loading