Skip to content

Commit 7d9b5f5

Browse files
FznamznonFznamznon
authored
[SYCL] Fix crash due to incorrect ReinterpretCastExpr generation (#7030)
When creating `ReinterpretCastExpr` use `getTrivialTypeSourceInfo` instead of `CreateTypeSourceInfo` since `CreateTypeSourceInfo` doesn't initialize underlying memory. The crash appeared during recursive AST visitors walks if kernel argument type had a nested name specifier.
1 parent a32021b commit 7d9b5f5

File tree

5 files changed

+189
-1
lines changed

5 files changed

+189
-1
lines changed

‎clang/lib/Sema/SemaSYCL.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2959,7 +2959,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
29592959
Expr *createReinterpretCastExpr(Expr *E, QualType To) {
29602960
return CXXReinterpretCastExpr::Create(
29612961
SemaRef.Context, To, VK_PRValue, CK_BitCast, E,
2962-
/*Path=*/nullptr, SemaRef.Context.CreateTypeSourceInfo(To),
2962+
/*Path=*/nullptr, SemaRef.Context.getTrivialTypeSourceInfo(To),
29632963
SourceLocation(), SourceLocation(), SourceRange());
29642964
}
29652965

‎clang/test/CodeGenSYCL/generated-types-initialization.cpp

Lines changed: 71 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,71 @@
1+
// RUN: %clang_cc1 -fno-sycl-force-inline-kernel-lambda -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -opaque-pointers -emit-llvm %s -o - | FileCheck %s
2+
3+
// This test checks that compiler generates correct code when kernel arguments
4+
// are structs that contain pointers but not decomposed.
5+
6+
#include "sycl.hpp"
7+
8+
struct A {
9+
float *F;
10+
};
11+
12+
struct B {
13+
int *F1;
14+
A F3;
15+
B(int *I, A AA) : F1(I), F3(AA) {};
16+
};
17+
18+
struct Nested {
19+
typedef B TDA;
20+
};
21+
22+
int main() {
23+
sycl::queue q;
24+
B Obj{nullptr, {nullptr}};
25+
26+
q.submit([&](sycl::handler &h) {
27+
h.single_task<class basic>(
28+
[=]() {
29+
(void)Obj;
30+
});
31+
});
32+
33+
Nested::TDA NNSObj{nullptr, {nullptr}};
34+
q.submit([&](sycl::handler &h) {
35+
h.single_task<class nns>([=]() {
36+
(void)NNSObj;
37+
});
38+
});
39+
return 0;
40+
}
41+
// CHECK: define dso_local spir_kernel void @{{.*}}basic(ptr noundef byval(%struct.__generated_B) align 8 %_arg_Obj)
42+
//
43+
// Kernel object clone.
44+
// CHECK: %[[K:[a-zA-Z0-9_.]+]] = alloca %class.anon
45+
// CHECK: %[[K_as_cast:[a-zA-Z0-9_.]+]] = addrspacecast ptr %[[K]] to ptr addrspace(4)
46+
//
47+
// Argument reference.
48+
// CHECK: %[[Arg_ref:[a-zA-Z0-9_.]+]] = addrspacecast ptr %_arg_Obj to ptr addrspace(4)
49+
//
50+
// Initialization.
51+
// CHECK: %[[GEP:[a-zA-Z0-9_.]+]] = getelementptr inbounds %class.anon, ptr addrspace(4) %[[K_as_cast]], i32 0, i32 0
52+
// CHECK: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 8 %[[GEP]], ptr addrspace(4) align 8 %[[Arg_ref]], i64 16, i1 false)
53+
//
54+
// Kernel body call.
55+
// CHECK: call spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_ENKUlvE_clEv(ptr addrspace(4) noundef align 8 dereferenceable_or_null(16) %[[K_as_cast]])
56+
57+
// CHECK: define dso_local spir_kernel void @{{.*}}nns(ptr noundef byval(%struct.__generated_B.0) align 8 %_arg_NNSObj)
58+
//
59+
// Kernel object clone.
60+
// CHECK: %[[NNSK:[a-zA-Z0-9_.]+]] = alloca %class.anon.2
61+
// CHECK: %[[NNSK_as_cast:[a-zA-Z0-9_.]+]] = addrspacecast ptr %[[NNSK]] to ptr addrspace(4)
62+
//
63+
// Argument reference.
64+
// CHECK: %[[NNSArg_ref:[a-zA-Z0-9_.]+]] = addrspacecast ptr %_arg_NNSObj to ptr addrspace(4)
65+
//
66+
// Initialization.
67+
// CHECK: %[[NNSGEP:[a-zA-Z0-9_.]+]] = getelementptr inbounds %class.anon.2, ptr addrspace(4) %[[NNSK_as_cast]], i32 0, i32 0
68+
// CHECK: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 8 %[[NNSGEP]], ptr addrspace(4) align 8 %[[NNSArg_ref]], i64 16, i1 false)
69+
//
70+
// Kernel body call.
71+
// CHECK: call spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_ENKUlvE_clEv(ptr addrspace(4) noundef align 8 dereferenceable_or_null(16) %[[NNSK_as_cast]])

‎clang/test/CodeGenSYCL/no-opaque-generated-types-initialization.cpp

Lines changed: 77 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,77 @@
1+
// RUN: %clang_cc1 -fno-sycl-force-inline-kernel-lambda -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -no-opaque-pointers -emit-llvm %s -o - | FileCheck %s
2+
3+
// This test checks that compiler generates correct code when kernel arguments
4+
// are structs that contain pointers but not decomposed.
5+
6+
#include "sycl.hpp"
7+
8+
struct A {
9+
float *F;
10+
};
11+
12+
struct B {
13+
int *F1;
14+
A F3;
15+
B(int *I, A AA) : F1(I), F3(AA) {};
16+
};
17+
18+
struct Nested {
19+
typedef B TDA;
20+
};
21+
22+
int main() {
23+
sycl::queue q;
24+
B Obj{nullptr, {nullptr}};
25+
26+
q.submit([&](sycl::handler &h) {
27+
h.single_task<class basic>(
28+
[=]() {
29+
(void)Obj;
30+
});
31+
});
32+
33+
Nested::TDA NNSObj{nullptr, {nullptr}};
34+
q.submit([&](sycl::handler &h) {
35+
h.single_task<class nns>([=]() {
36+
(void)NNSObj;
37+
});
38+
});
39+
return 0;
40+
}
41+
// CHECK: define dso_local spir_kernel void @{{.*}}basic(%struct.__generated_B* noundef byval(%struct.__generated_B) align 8 %_arg_Obj)
42+
//
43+
// Kernel object clone.
44+
// CHECK: %[[K:[a-zA-Z0-9_.]+]] = alloca %class.anon
45+
// CHECK: %[[K_as_cast:[a-zA-Z0-9_.]+]] = addrspacecast %class.anon* %[[K]] to %class.anon addrspace(4)*
46+
//
47+
// Argument reference.
48+
// CHECK: %[[Arg_ref:[a-zA-Z0-9_.]+]] = addrspacecast %struct.__generated_B* %_arg_Obj to %struct.__generated_B addrspace(4)*
49+
50+
// Initialization.
51+
// CHECK: %[[GEP:[a-zA-Z0-9_.]+]] = getelementptr inbounds %class.anon, %class.anon addrspace(4)* %[[K_as_cast]], i32 0, i32 0
52+
// CHECK: %[[ArgBC:[a-zA-Z0-9_.]+]] = bitcast %struct.__generated_B addrspace(4)* %[[Arg_ref]] to %struct.B addrspace(4)*
53+
// CHECK: %[[GEPBC:[a-zA-Z0-9_.]+]] = bitcast %struct.B addrspace(4)* %[[GEP]] to i8 addrspace(4)*
54+
// CHECK: %[[ArgBC2:[a-zA-Z0-9_.]+]] = bitcast %struct.B addrspace(4)* %[[ArgBC]] to i8 addrspace(4)*
55+
// CHECK: call void @llvm.memcpy.p4i8.p4i8.i64(i8 addrspace(4)* align 8 %[[GEPBC]], i8 addrspace(4)* align 8 %[[ArgBC2]], i64 16, i1 false)
56+
//
57+
// Kernel body call.
58+
// CHECK: call spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_ENKUlvE_clEv(%class.anon addrspace(4)* noundef align 8 dereferenceable_or_null(16) %[[K_as_cast]])
59+
60+
// CHECK: define dso_local spir_kernel void @{{.*}}nns(%struct.__generated_B.0* noundef byval(%struct.__generated_B.0) align 8 %_arg_NNSObj)
61+
//
62+
// Kernel object clone.
63+
// CHECK: %[[NNSK:[a-zA-Z0-9_.]+]] = alloca %class.anon.2
64+
// CHECK: %[[NNSK_as_cast:[a-zA-Z0-9_.]+]] = addrspacecast %class.anon.2* %[[NNSK]] to %class.anon.2 addrspace(4)*
65+
//
66+
// Argument reference.
67+
// CHECK: %[[NNSArg_ref:[a-zA-Z0-9_.]+]] = addrspacecast %struct.__generated_B.0* %_arg_NNSObj to %struct.__generated_B.0 addrspace(4)*
68+
//
69+
// Initialization.
70+
// CHECK: %[[NNSGEP:[a-zA-Z0-9_.]+]] = getelementptr inbounds %class.anon.2, %class.anon.2 addrspace(4)* %[[NNSK_as_cast]], i32 0, i32 0
71+
// CHECK: %[[NNSArgBC:[a-zA-Z0-9_.]+]] = bitcast %struct.__generated_B.0 addrspace(4)* %[[NNSArg_ref]] to %struct.B addrspace(4)*
72+
// CHECK: %[[NNSGEPBC:[a-zA-Z0-9_.]+]] = bitcast %struct.B addrspace(4)* %[[NNSGEP]] to i8 addrspace(4)*
73+
// CHECK: %[[NNSArgBC2:[a-zA-Z0-9_.]+]] = bitcast %struct.B addrspace(4)* %[[NNSArgBC]] to i8 addrspace(4)*
74+
// CHECK: call void @llvm.memcpy.p4i8.p4i8.i64(i8 addrspace(4)* align 8 %[[NNSGEPBC]], i8 addrspace(4)* align 8 %[[NNSArgBC2]], i64 16, i1 false)
75+
//
76+
// Kernel body call.
77+
// CHECK: call spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_ENKUlvE_clEv(%class.anon.2 addrspace(4)* noundef align 8 dereferenceable_or_null(16) %[[NNSK_as_cast]])

‎clang/test/SemaSYCL/built-in-type-kernel-arg.cpp

Lines changed: 30 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -23,6 +23,15 @@ void test(const int some_const) {
2323
});
2424
}
2525

26+
struct test_struct_simple {
27+
int data;
28+
int *ptr;
29+
};
30+
31+
struct Nested {
32+
typedef test_struct_simple TDS;
33+
};
34+
2635
int main() {
2736
int data = 5;
2837
int* data_addr = &data;
@@ -54,6 +63,15 @@ int main() {
5463
});
5564
});
5665

66+
Nested::TDS tds;
67+
deviceQueue.submit([&](sycl::handler &h) {
68+
h.single_task<class kernel_nns>(
69+
[=]() {
70+
test_struct_simple k_s;
71+
k_s = tds;
72+
});
73+
});
74+
5775
const int some_const = 10;
5876
test(some_const);
5977
return 0;
@@ -162,3 +180,15 @@ int main() {
162180
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' <AddressSpaceConversion>
163181
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' <LValueToRValue>
164182
// CHECK-NEXT: DeclRefExpr {{.*}} '__global int *' lvalue ParmVar {{.*}} '_arg_ptr_array' '__global int *'
183+
184+
// CHECK: FunctionDecl {{.*}}kernel_nns 'void (__generated_test_struct_simple)'
185+
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_tds '__generated_test_struct_simple'
186+
187+
// CHECK: VarDecl {{.*}} used __SYCLKernel
188+
// CHECK: InitListExpr
189+
// CHECK: CXXConstructExpr {{.*}} 'Nested::TDS':'test_struct_simple' 'void (const test_struct_simple &) noexcept'
190+
// CHECK: ImplicitCastExpr {{.*}} 'const test_struct_simple' lvalue <NoOp>
191+
// CHECK: UnaryOperator {{.*}} 'Nested::TDS':'test_struct_simple' lvalue prefix '*' cannot overflow
192+
// CHECK: CXXReinterpretCastExpr {{.*}} 'Nested::TDS *' reinterpret_cast<struct Nested::TDS *> <BitCast>
193+
// CHECK: UnaryOperator {{.*}} '__generated_test_struct_simple *' prefix '&' cannot overflow
194+
// CHECK: DeclRefExpr {{.*}} '__generated_test_struct_simple' lvalue ParmVar {{.*}} '_arg_tds' '__generated_test_struct_simple'

‎clang/test/SemaSYCL/decomposition.cpp

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -53,6 +53,10 @@ struct StructWithPtr {
5353
int i;
5454
};
5555

56+
struct Nested {
57+
typedef StructWithPtr TDStrWithPTR;
58+
};
59+
5660
struct NonTrivialType {
5761
int *Ptr;
5862
int i;
@@ -179,6 +183,12 @@ int main() {
179183
});
180184
// CHECK: FunctionDecl {{.*}}Pointer{{.*}} 'void (__generated_StructWithPtr)'
181185

186+
Nested::TDStrWithPTR TDStructWithPtr;
187+
myQueue.submit([&](sycl::handler &h) {
188+
h.single_task<class TDStr>([=]() { return TDStructWithPtr.i; });
189+
});
190+
// CHECK: FunctionDecl {{.*}}TDStr{{.*}} 'void (__generated_StructWithPtr)'
191+
182192
// FIXME: Stop decomposition of arrays with pointers
183193
StructWithArray<StructWithPtr> t1;
184194
myQueue.submit([&](sycl::handler &h) {

0 commit comments

Comments
 (0)