Skip to content

Commit e56db5f

Browse files
zhaomaosuzhaomaosu
committed
Merge branch 'sycl' into san-spec-constant-devicety
2 parents 71a5527 + 84518c1 commit e56db5f

File tree

93 files changed

+1099
-359
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

93 files changed

+1099
-359
lines changed

‎clang/lib/AST/ASTContext.cpp‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -12824,7 +12824,8 @@ static GVALinkage adjustGVALinkageForAttributes(const ASTContext &Context,
1282412824
if (Context.shouldExternalize(D))
1282512825
return GVA_StrongExternal;
1282612826
} else if (Context.getLangOpts().SYCLIsDevice &&
12827-
D->hasAttr<DeviceKernelAttr>()) {
12827+
(D->hasAttr<DeviceKernelAttr>() &&
12828+
D->getAttr<DeviceKernelAttr>()->isImplicit())) {
1282812829
if (L == GVA_DiscardableODR)
1282912830
return GVA_StrongODR;
1283012831
}

‎clang/lib/Sema/SemaSYCL.cpp‎

Lines changed: 14 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -4097,10 +4097,13 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
40974097
}
40984098

40994099
void addArrayElementInit(FieldDecl *FD, QualType T) {
4100-
Expr *RCE = createReinterpretCastExpr(
4101-
createGetAddressOf(ArrayParamBases.pop_back_val()),
4102-
SemaSYCLRef.getASTContext().getPointerType(T));
4103-
Expr *Initializer = createDerefOp(RCE);
4100+
Expr *Initializer = ArrayParamBases.pop_back_val();
4101+
if (!T->isPointerType()) {
4102+
Expr *RCE = createReinterpretCastExpr(
4103+
createGetAddressOf(Initializer),
4104+
SemaSYCLRef.getASTContext().getPointerType(T));
4105+
Initializer = createDerefOp(RCE);
4106+
}
41044107
addFieldInit(FD, T, Initializer);
41054108
}
41064109

@@ -5443,9 +5446,13 @@ void SemaSYCL::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc,
54435446
ESIMDKernelDiagnostics esimdKernel(*this, KernelObj->getLocation(),
54445447
IsSIMDKernel);
54455448

5446-
SyclKernelDeclCreator kernel_decl(*this, KernelObj->getLocation(),
5447-
KernelCallerFunc->isInlined(), IsSIMDKernel,
5448-
KernelCallerFunc);
5449+
// In case of syntax errors in input programs we are not able to access
5450+
// CallOperator. In this case the value of IsInlined doesn't matter, because
5451+
// compilation will fail with errors anyways.
5452+
const bool IsInlined =
5453+
CallOperator ? CallOperator->isInlined() : /* placeholder */ false;
5454+
SyclKernelDeclCreator kernel_decl(*this, KernelObj->getLocation(), IsInlined,
5455+
IsSIMDKernel, KernelCallerFunc);
54495456
SyclKernelBodyCreator kernel_body(*this, kernel_decl, KernelObj,
54505457
KernelCallerFunc, IsSIMDKernel,
54515458
CallOperator);

‎clang/test/CodeGenSYCL/array-kernel-param.cpp‎

Lines changed: 66 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,66 @@
1+
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5
2+
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
3+
4+
#include "sycl.hpp"
5+
6+
sycl::queue myQueue;
7+
8+
using namespace sycl;
9+
10+
// CHECK: %struct.__wrapper_class = type { [2 x i32] }
11+
// CHECK: %class.anon = type { [2 x i32] }
12+
// CHECK: %struct.__wrapper_class.0 = type { [2 x ptr addrspace(1)] }
13+
// CHECK: %class.anon.1 = type { [2 x ptr addrspace(4)] }
14+
15+
int main() {
16+
int Array[2];
17+
myQueue.submit([&](sycl::handler &h) {
18+
h.single_task<class IntArray>(
19+
[=] {
20+
int local = Array[1];
21+
});
22+
});
23+
24+
// CHECK-LABEL: @{{.*}}IntArray(ptr {{.*}}byval(%struct.__wrapper_class)
25+
// CHECK: %__SYCLKernel = alloca %class.anon, align 4
26+
// CHECK: %__SYCLKernel.ascast = addrspacecast ptr %__SYCLKernel to ptr addrspace(4)
27+
// CHECK: %_arg_Array.ascast = addrspacecast ptr %_arg_Array to ptr addrspace(4)
28+
// CHECK: %Array = getelementptr inbounds nuw %class.anon, ptr addrspace(4) %__SYCLKernel.ascast, i32 0, i32 0
29+
// CHECK: %0 = getelementptr inbounds nuw %struct.__wrapper_class, ptr addrspace(4) %_arg_Array.ascast, i32 0, i32 0
30+
// CHECK: %arrayinit.begin = getelementptr inbounds [2 x i32], ptr addrspace(4) %Array, i64 0, i64 0
31+
// CHECK: br label %arrayinit.body
32+
// CHECK: arrayinit.body: ; preds = %arrayinit.body, %entry
33+
// CHECK: %arrayinit.index = phi i64 [ 0, %entry ], [ %arrayinit.next, %arrayinit.body ]
34+
// CHECK: %1 = getelementptr inbounds i32, ptr addrspace(4) %arrayinit.begin, i64 %arrayinit.index
35+
// CHECK: %arrayidx = getelementptr inbounds nuw [2 x i32], ptr addrspace(4) %0, i64 0, i64 %arrayinit.index
36+
// CHECK: %2 = load i32, ptr addrspace(4) %arrayidx, align 4
37+
// CHECK: store i32 %2, ptr addrspace(4) %1, align 4
38+
// CHECK: %arrayinit.next = add nuw i64 %arrayinit.index, 1
39+
// CHECK: %arrayinit.done = icmp eq i64 %arrayinit.next, 2
40+
// CHECK: br i1 %arrayinit.done, label %arrayinit.end, label %arrayinit.body
41+
// CHECK: arrayinit.end: ; preds = %arrayinit.body
42+
43+
int *ArrayOfPointers[2];
44+
myQueue.submit([&](sycl::handler &h) {
45+
h.single_task<class PtrArray>(
46+
[=] {
47+
int local = *ArrayOfPointers[1];
48+
});
49+
});
50+
// CHECK-LABEL: @{{.*}}PtrArray(ptr {{.*}}byval(%struct.__wrapper_class.0)
51+
// CHECK: %__SYCLKernel = alloca %class.anon.1, align 8
52+
// CHECK: %__SYCLKernel.ascast = addrspacecast ptr %__SYCLKernel to ptr addrspace(4)
53+
// CHECK: %_arg_ArrayOfPointers.ascast = addrspacecast ptr %_arg_ArrayOfPointers to ptr addrspace(4)
54+
// CHECK: %ArrayOfPointers = getelementptr inbounds nuw %class.anon.1, ptr addrspace(4) %__SYCLKernel.ascast, i32 0, i32 0
55+
// CHECK: %0 = getelementptr inbounds nuw %struct.__wrapper_class.0, ptr addrspace(4) %_arg_ArrayOfPointers.ascast, i32 0, i32 0
56+
// CHECK: %arrayidx = getelementptr inbounds nuw [2 x ptr addrspace(1)], ptr addrspace(4) %0, i64 0, i64 0
57+
// CHECK: %1 = load ptr addrspace(1), ptr addrspace(4) %arrayidx, align 8
58+
// CHECK: %2 = addrspacecast ptr addrspace(1) %1 to ptr addrspace(4)
59+
// CHECK: store ptr addrspace(4) %2, ptr addrspace(4) %ArrayOfPointers, align 8
60+
// CHECK: %arrayinit.element = getelementptr inbounds ptr addrspace(4), ptr addrspace(4) %ArrayOfPointers, i64 1
61+
// CHECK: %3 = getelementptr inbounds nuw %struct.__wrapper_class.0, ptr addrspace(4) %_arg_ArrayOfPointers.ascast, i32 0, i32 0
62+
// CHECK: %arrayidx1 = getelementptr inbounds nuw [2 x ptr addrspace(1)], ptr addrspace(4) %3, i64 0, i64 1
63+
// CHECK: %4 = load ptr addrspace(1), ptr addrspace(4) %arrayidx1, align 8
64+
// CHECK: %5 = addrspacecast ptr addrspace(1) %4 to ptr addrspace(4)
65+
// CHECK: store ptr addrspace(4) %5, ptr addrspace(4) %arrayinit.element, align 8
66+
}

‎clang/test/CodeGenSYCL/device_has.cpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -6,7 +6,7 @@
66
using namespace sycl;
77
queue q;
88

9-
// CHECK-DAG: define dso_local spir_kernel void @{{.*}}kernel_name_1{{.*}} !sycl_declared_aspects ![[ASPECTS1:[0-9]+]] !srcloc ![[SRCLOC1:[0-9]+]]
9+
// CHECK-DAG: define {{.*}}spir_kernel void @{{.*}}kernel_name_1{{.*}} !sycl_declared_aspects ![[ASPECTS1:[0-9]+]] !srcloc ![[SRCLOC1:[0-9]+]]
1010

1111
// CHECK-DAG: define {{.*}}spir_func void @{{.*}}func1{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC2:[0-9]+]] {
1212
[[sycl::device_has(sycl::aspect::cpu)]] void func1() {}
@@ -67,7 +67,7 @@ void foo() {
6767
q.submit([&](handler &h) {
6868
KernelFunctor f1;
6969
h.single_task<class kernel_name_1>(f1);
70-
// CHECK-DAG: define dso_local spir_kernel void @{{.*}}kernel_name_2{{.*}} !sycl_declared_aspects ![[ASPECTS4:[0-9]+]] !srcloc ![[SRCLOC8:[0-9]+]]
70+
// CHECK-DAG: define {{.*}}spir_kernel void @{{.*}}kernel_name_2{{.*}} !sycl_declared_aspects ![[ASPECTS4:[0-9]+]] !srcloc ![[SRCLOC8:[0-9]+]]
7171
h.single_task<class kernel_name_2>([]() [[sycl::device_has(sycl::aspect::gpu)]] {});
7272
});
7373
}

‎clang/test/CodeGenSYCL/dynamic_local_accessor.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -7,7 +7,7 @@
77
// The first two RUN commands verify that the init call is generated with the correct arguments in LLVM IR
88
// and the second two RUN commands verify the contents of the integration header produced by the frontend.
99
//
10-
// CHECK-IR: define dso_local spir_kernel void @
10+
// CHECK-IR: define {{.*}}spir_kernel void @
1111
// CHECK-IR-SAME: ptr addrspace(3) noundef align 4 [[PTR:%[a-zA-Z0-9_]+]]
1212
//
1313
// CHECK-IR: [[PTR]].addr = alloca ptr addrspace(3), align 8

‎clang/test/CodeGenSYCL/dynamic_work_group_memory.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -7,7 +7,7 @@
77
// The first two RUN commands verify that the init call is generated with the correct arguments in LLVM IR
88
// and the second two RUN commands verify the contents of the integration header produced by the frontend.
99
//
10-
// CHECK-IR: define dso_local spir_kernel void @
10+
// CHECK-IR: define {{.*}}spir_kernel void @
1111
// CHECK-IR-SAME: ptr addrspace(3) noundef align 4 [[PTR:%[a-zA-Z0-9_]+]]
1212
//
1313
// CHECK-IR: [[PTR]].addr = alloca ptr addrspace(3), align 8

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -38,7 +38,7 @@ int main() {
3838
});
3939
return 0;
4040
}
41-
// CHECK: define dso_local spir_kernel void @{{.*}}basic(ptr noundef byval(%struct.__generated_B) align 8 %_arg_Obj)
41+
// CHECK: define {{.*}}spir_kernel void @{{.*}}basic(ptr noundef byval(%struct.__generated_B) align 8 %_arg_Obj)
4242
//
4343
// Kernel object clone.
4444
// CHECK: %[[K:[a-zA-Z0-9_.]+]] = alloca %class.anon
@@ -54,7 +54,7 @@ int main() {
5454
// Kernel body call.
5555
// CHECK: call spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_ENKUlvE_clEv(ptr addrspace(4) noundef align 8 dereferenceable_or_null(16) %[[K_as_cast]])
5656

57-
// CHECK: define dso_local spir_kernel void @{{.*}}nns(ptr noundef byval(%struct.__generated_B.0) align 8 %_arg_NNSObj)
57+
// CHECK: define {{.*}}spir_kernel void @{{.*}}nns(ptr noundef byval(%struct.__generated_B.0) align 8 %_arg_NNSObj)
5858
//
5959
// Kernel object clone.
6060
// CHECK: %[[NNSK:[a-zA-Z0-9_.]+]] = alloca %class.anon.2

‎clang/test/CodeGenSYCL/kernel-handler.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -22,7 +22,7 @@ void test(int val) {
2222
});
2323
}
2424

25-
// ALL: define dso_local{{ spir_kernel | ptx_kernel | }}void @{{.*}}test_kernel_handler{{[^(]*}}
25+
// ALL: define {{.*}}{{ spir_kernel | ptx_kernel | }}void @{{.*}}test_kernel_handler{{[^(]*}}
2626
// NONATIVESUPPORT-SAME: (ptr noundef byval(%class.anon) align 4 %_arg__sycl_functor, ptr addrspace(1) noundef align 1 %_arg__specialization_constants_buffer)
2727
// NATIVESUPPORT-SAME: (i32 noundef %_arg_a, ptr addrspace(1) noundef align 1 %_arg__specialization_constants_buffer)
2828
// ALL: %kh = alloca %"class.sycl::_V1::kernel_handler", align 1

‎clang/test/CodeGenSYCL/kernel-op-calls.cpp‎

Lines changed: 11 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -14,46 +14,46 @@ class Functor1 {
1414
[[sycl::reqd_sub_group_size(4)]] void operator()(sycl::id<1> id) const {}
1515

1616
[[sycl::work_group_size_hint(1, 2, 3)]] void operator()(sycl::id<2> id) const {}
17-
1817
};
1918

2019
class ESIMDFunctor {
2120
public:
22-
ESIMDFunctor(){}
21+
ESIMDFunctor(){}
2322

2423
[[intel::sycl_explicit_simd]] void operator()(sycl::id<2> id) const {}
2524

26-
[[sycl::work_group_size_hint(1, 2, 3)]][[intel::sycl_explicit_simd]] void operator()(sycl::id<1> id) const {}
27-
25+
[[sycl::work_group_size_hint(1, 2, 3)]] [[intel::sycl_explicit_simd]]
26+
void operator()(sycl::id<1> id) const {}
2827
};
2928

3029
// Check templated 'operator()()' call works.
3130
class kernels {
32-
public:
31+
public:
3332
kernels(){}
3433

35-
template<int Dimensions = 1>
36-
[[sycl::work_group_size_hint(1, 2, 3)]] void operator()(sycl::id<Dimensions> item) const {}
37-
34+
template<int Dimensions = 1>
35+
[[sycl::work_group_size_hint(1, 2, 3)]]
36+
void operator()(sycl::id<Dimensions> item) const {}
3837
};
3938

4039
int main() {
4140

4241
Q.submit([&](sycl::handler& cgh) {
4342
Functor1 F;
44-
// CHECK: define dso_local spir_kernel void @_ZTS8Functor1() {{.*}} !kernel_arg_buffer_location !{{[0-9]+}} !intel_reqd_sub_group_size !{{[0-9]+}} !sycl_fixed_targets !{{[0-9]+}} {
43+
// CHECK: define weak_odr spir_kernel void @_ZTS8Functor1() {{.*}} !intel_reqd_sub_group_size
4544
cgh.parallel_for(sycl::range<1>(10), F);
4645
});
4746

4847
Q.submit([&](sycl::handler& cgh) {
4948
kernels K;
50-
// CHECK: define dso_local spir_kernel void @_ZTS7kernels() {{.*}} !kernel_arg_buffer_location !{{[0-9]+}} !work_group_size_hint !{{[0-9]+}} !sycl_fixed_targets !{{[0-9]+}} {
49+
// CHECK: define weak_odr spir_kernel void @_ZTS7kernels() {{.*}} !work_group_size_hint !{{[0-9]+}}
5150
cgh.parallel_for(sycl::range<1>(10), K);
5251
});
5352

5453
Q.submit([&](sycl::handler& cgh) {
5554
ESIMDFunctor EF;
56-
// CHECK: define dso_local spir_kernel void @_ZTS12ESIMDFunctor() {{.*}} !sycl_explicit_simd !{{[0-9]+}} !intel_reqd_sub_group_size !{{[0-9]+}} !work_group_size_hint !{{[0-9]+}} !kernel_arg_accessor_ptr !{{[0-9]+}} !sycl_fixed_targets !{{[0-9]+}} {
55+
// CHECK: define weak_odr spir_kernel void @_ZTS12ESIMDFunctor() {{.*}} !work_group_size_hint
56+
// CHECK-SAME: !sycl_explicit_simd
5757
cgh.parallel_for(sycl::range<1>(10), EF);
5858
});
5959

‎clang/test/CodeGenSYCL/kernel-param-acc-array.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -24,7 +24,7 @@ int main() {
2424
acc[1].use();
2525
});
2626
}
27-
// CHECK-LABEL: define dso_local spir_kernel void @_ZTSZ4mainE8kernel_A(
27+
// CHECK-LABEL: define {{.*}}spir_kernel void @_ZTSZ4mainE8kernel_A(
2828
// CHECK-SAME: ptr addrspace(1) noundef align 4 [[_ARG_ACC:%.*]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[_ARG_ACC1:%.*]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[_ARG_ACC2:%.*]], ptr noundef byval(%"struct.sycl::_V1::id") align 4 [[_ARG_ACC3:%.*]], ptr addrspace(1) noundef align 4 [[_ARG_ACC4:%.*]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[_ARG_ACC6:%.*]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[_ARG_ACC7:%.*]], ptr noundef byval(%"struct.sycl::_V1::id") align 4 [[_ARG_ACC8:%.*]]) #[[ATTR0:[0-9]+]]
2929
// CHECK-NEXT: [[ENTRY:.*:]]
3030
// CHECK-NEXT: [[_ARG_ACC_ADDR:%.*]] = alloca ptr addrspace(1), align 8

0 commit comments

Comments
 (0)