Skip to content

Commit e6e086b

Browse files
author
iclsrc
committed
Merge from 'sycl' to 'sycl-web'
2 parents 5108ef8 + cf0d053 commit e6e086b

File tree

157 files changed

+605
-599
lines changed

Some content is hidden

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

157 files changed

+605
-599
lines changed

‎clang/test/CodeGenSYCL/accessor_inheritance.cpp

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -67,13 +67,13 @@ int main() {
6767
// CHECK: [[ACC_FIELD:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base* [[BITCAST]], i32 0, i32 2
6868
// CHECK: [[ACC1_AS_CAST:%[a-zA-Z0-9_]+]] = addrspacecast %"class{{.*}}cl::sycl::accessor"* [[ACC_FIELD]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)*
6969
// Default constructor call
70-
// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEC1Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[ACC1_AS_CAST]])
70+
// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEC1Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[ACC1_AS_CAST]])
7171
// CHECK: [[BITCAST1:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured* [[GEP]] to i8*
7272
// CHECK: [[GEP1:%[a-zA-Z0-9_]+]] = getelementptr inbounds i8, i8* [[BITCAST1]], i64 20
7373
// CHECK: [[BITCAST2:%[a-zA-Z0-9_]+]] = bitcast i8* [[GEP1]] to %"class{{.*}}cl::sycl::accessor"*
7474
// CHECK: [[ACC2_AS_CAST:%[a-zA-Z0-9_]+]] = addrspacecast %"class{{.*}}cl::sycl::accessor"* [[BITCAST2]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)*
7575
// Default constructor call
76-
// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEC2Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[ACC2_AS_CAST]])
76+
// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEC2Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[ACC2_AS_CAST]])
7777

7878
// CHECK C field initialization
7979
// CHECK: [[FIELD_C:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Captured, %struct{{.*}}Captured* [[GEP]], i32 0, i32 2
@@ -86,9 +86,9 @@ int main() {
8686
// CHECK: [[ACC1_FIELD:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base* [[BITCAST3]], i32 0, i32 2
8787
// CHECK: [[ACC1_DATA_LOAD:%[a-zA-Z0-9_]+]] = load i8 addrspace(1)*, i8 addrspace(1)** [[ACC1_DATA]].addr
8888
// CHECK: [[ACC1_AS_CAST1:%[a-zA-Z0-9_]+]] = addrspacecast %"class{{.*}}cl::sycl::accessor"* [[ACC1_FIELD]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)*
89-
// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[ACC1_AS_CAST1]], i8 addrspace(1)* [[ACC1_DATA_LOAD]]
89+
// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class{{.*}}cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[ACC1_AS_CAST1]], i8 addrspace(1)* [[ACC1_DATA_LOAD]]
9090
//
9191
// CHECK: [[GEP3:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class._ZTSZ4mainE3$_0.anon", %"class._ZTSZ4mainE3$_0.anon"* [[KERNEL_OBJ]], i32 0, i32 0
9292
// CHECK: [[ACC2_DATA_LOAD:%[a-zA-Z0-9_]+]] = load i8 addrspace(1)*, i8 addrspace(1)** [[ACC2_DATA]].addr
9393
// CHECK: [[AS_CAST_CAPTURED:%[a-zA-Z0-9_]+]] = addrspacecast %struct{{.*}}Captured* [[GEP3]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)*
94-
// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[AS_CAST_CAPTURED]], i8 addrspace(1)* [[ACC2_DATA_LOAD]]
94+
// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class{{.*}}cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[AS_CAST_CAPTURED]], i8 addrspace(1)* [[ACC2_DATA_LOAD]]

‎clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -48,9 +48,9 @@ int main() {
4848
// Check accessor __init method call
4949
// CHECK-OLD: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor"* [[ACCESSOR]], i32 addrspace(1)* [[MEM_LOAD]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ARANGE]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MRANGE]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OID]])
5050
// CHECK: [[ACCESSORCAST:%[0-9]+]] = addrspacecast %"class{{.*}}accessor"* [[ACCESSOR]] to %"class{{.*}}accessor" addrspace(4)*
51-
// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* [[ACCESSORCAST]], i32 addrspace(1)* [[MEM_LOAD]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ARANGE]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MRANGE]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OID]])
51+
// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[ACCESSORCAST]], i32 addrspace(1)* [[MEM_LOAD]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ARANGE]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MRANGE]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OID]])
5252

5353
// Check lambda "()" operator call
5454
// CHECK-OLD: call spir_func void @{{.*}}(%"class.{{.*}}.anon"* [[ANON]])
5555
// CHECK: [[ANONCAST:%[0-9]+]] = addrspacecast %"class{{.*}}anon"* {{.*}} to %"class{{.*}}anon" addrspace(4)*
56-
// CHECK: call spir_func void @{{.*}}(%"class.{{.*}}.anon" addrspace(4)* [[ANONCAST]])
56+
// CHECK: call spir_func void @{{.*}}(%"class.{{.*}}.anon" addrspace(4)* {{[^,]*}} [[ANONCAST]])

‎clang/test/CodeGenSYCL/device-functions.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -22,6 +22,6 @@ int main() {
2222
return 0;
2323
}
2424
// CHECK: define spir_kernel void @_ZTSZ4mainE11fake_kernel()
25-
// CHECK: define internal spir_func void @"_ZZ4mainENK3$_0clEv"(%"class.{{.*}}.anon" addrspace(4)* %this)
25+
// CHECK: define internal spir_func void @"_ZZ4mainENK3$_0clEv"(%"class.{{.*}}.anon" addrspace(4)* {{[^,]*}} %this)
2626
// CHECK: define spir_func void @_Z3foov()
2727
// CHECK: define linkonce_odr spir_func i32 @_Z3barIiET_S0_(i32 %arg)

‎clang/test/CodeGenSYCL/esimd-accessor-ptr-md.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -38,7 +38,7 @@ void test(int val) {
3838
// --- Attributes
3939
// CHECK: {{.*}} !kernel_arg_accessor_ptr ![[ACC_PTR_ATTR:[0-9]+]] !sycl_explicit_simd !{{[0-9]+}} {{.*}}{
4040
// --- init_esimd call is expected instead of __init:
41-
// CHECK: call spir_func void @{{.*}}__init_esimd{{.*}}(%"{{.*}}sycl::accessor" addrspace(4)* %{{[0-9]+}}, i32 addrspace(1)* %{{[0-9]+}})
41+
// CHECK: call spir_func void @{{.*}}__init_esimd{{.*}}(%"{{.*}}sycl::accessor" addrspace(4)* {{[^,]*}} %{{[0-9]+}}, i32 addrspace(1)* %{{[0-9]+}})
4242
// CHECK-LABEL: }
4343
// CHECK: ![[ACC_PTR_ATTR]] = !{i1 true, i1 false, i1 true}
4444
}

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

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -55,24 +55,24 @@ int main() {
5555
// Clang takes advantage of element 1 having the same address as the array, so it doesn't do a GEP.
5656
// CHECK: [[ELEM1_ASCAST:%[a-zA-Z0-9_.]+]] = addrspacecast [[ACCESSOR]]* [[BEGIN]] to [[ACCESSOR]] addrspace(4)*
5757
// CTOR Call #1
58-
// CHECK: call spir_func void @{{.+}}([[ACCESSOR]] addrspace(4)* [[ELEM1_ASCAST]])
58+
// CHECK: call spir_func void @{{.+}}([[ACCESSOR]] addrspace(4)* {{[^,]*}} [[ELEM1_ASCAST]])
5959
// CHECK: [[ELEM2_GEP:%[a-zA-Z0-9_.]+]] = getelementptr inbounds [[ACCESSOR]], [[ACCESSOR]]* [[BEGIN]], i64 1
6060
// CHECK: [[ELEM2_ASCAST:%[a-zA-Z0-9_.]+]] = addrspacecast [[ACCESSOR]]* [[ELEM2_GEP]] to [[ACCESSOR]] addrspace(4)*
6161
// CTOR Call #2
62-
// CHECK: call spir_func void @{{.+}}([[ACCESSOR]] addrspace(4)* [[ELEM2_ASCAST]])
62+
// CHECK: call spir_func void @{{.+}}([[ACCESSOR]] addrspace(4)* {{[^,]*}} [[ELEM2_ASCAST]])
6363

6464
// CHECK acc[0] __init method call
6565
// CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
6666
// CHECK: [[INDEX1:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], [2 x [[ACCESSOR]]]* [[ACCESSOR_ARRAY1]], i64 0, i64 0
6767
// CHECK load from kernel pointer argument alloca
6868
// CHECK: [[MEM_LOAD1:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG1]]
6969
// CHECK: [[ACC_CAST1:%[0-9]+]] = addrspacecast [[ACCESSOR]]* [[INDEX1]] to [[ACCESSOR]] addrspace(4)*
70-
// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* [[ACC_CAST1]], i32 addrspace(1)* [[MEM_LOAD1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE1]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET1]])
70+
// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[ACC_CAST1]], i32 addrspace(1)* [[MEM_LOAD1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE1]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET1]])
7171

7272
// CHECK acc[1] __init method call
7373
// CHECK: [[ACCESSOR_ARRAY2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
7474
// CHECK: [[INDEX2:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], [2 x [[ACCESSOR]]]* [[ACCESSOR_ARRAY2]], i64 0, i64 1
7575
// CHECK load from kernel pointer argument alloca
7676
// CHECK: [[MEM_LOAD2:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG2]]
7777
// CHECK: [[ACC_CAST2:%[0-9]+]] = addrspacecast [[ACCESSOR]]* [[INDEX2]] to [[ACCESSOR]] addrspace(4)*
78-
// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* [[ACC_CAST2]], i32 addrspace(1)* [[MEM_LOAD2]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE2]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE2]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET2]])
78+
// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[ACC_CAST2]], i32 addrspace(1)* [[MEM_LOAD2]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE2]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE2]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET2]])

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

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -59,24 +59,24 @@ int main() {
5959
// Clang takes advantage of element 1 having the same address as the array, so it doesn't do a GEP.
6060
// CHECK: [[ELEM1_ASCAST:%[a-zA-Z0-9_.]+]] = addrspacecast [[ACCESSOR]]* [[BEGIN]] to [[ACCESSOR]] addrspace(4)*
6161
// CTOR Call #1
62-
// CHECK: call spir_func void @{{.+}}([[ACCESSOR]] addrspace(4)* [[ELEM1_ASCAST]])
62+
// CHECK: call spir_func void @{{.+}}([[ACCESSOR]] addrspace(4)* {{[^,]*}} [[ELEM1_ASCAST]])
6363
// CHECK: [[ELEM2_GEP:%[a-zA-Z0-9_.]+]] = getelementptr inbounds [[ACCESSOR]], [[ACCESSOR]]* [[BEGIN]], i64 1
6464
// CHECK: [[ELEM2_ASCAST:%[a-zA-Z0-9_.]+]] = addrspacecast [[ACCESSOR]]* [[ELEM2_GEP]] to [[ACCESSOR]] addrspace(4)*
6565
// CTOR Call #2
66-
// CHECK: call spir_func void @{{.+}}([[ACCESSOR]] addrspace(4)* [[ELEM2_ASCAST]])
66+
// CHECK: call spir_func void @{{.+}}([[ACCESSOR]] addrspace(4)* {{[^,]*}} [[ELEM2_ASCAST]])
6767

6868
// Check acc[0] __init method call
6969
// CHECK: [[GEP_LAMBDA1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
7070
// CHECK: [[GEP_MEMBER_ACC1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}.struct_acc_t, %struct.{{.*}}.struct_acc_t* [[GEP_LAMBDA1]], i32 0, i32 0
7171
// CHECK: [[ARRAY_IDX1:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], [2 x [[ACCESSOR]]]* [[GEP_MEMBER_ACC1]], i64 0, i64 0
7272
// CHECK: [[MEM_LOAD1:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG1]].addr
7373
// CHECK: [[ACC_CAST1:%[0-9]+]] = addrspacecast [[ACCESSOR]]* [[ARRAY_IDX1]] to [[ACCESSOR]] addrspace(4)*
74-
// CHECK: call spir_func void @{{.*}}__init{{.*}}([[ACCESSOR]] addrspace(4)* [[ACC_CAST1]], i32 addrspace(1)* [[MEM_LOAD1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE1]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET1]])
74+
// CHECK: call spir_func void @{{.*}}__init{{.*}}([[ACCESSOR]] addrspace(4)* {{[^,]*}} [[ACC_CAST1]], i32 addrspace(1)* [[MEM_LOAD1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE1]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET1]])
7575

7676
// Check acc[1] __init method call
7777
// CHECK: [[GEP_LAMBDA2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
7878
// CHECK: [[GEP_MEMBER_ACC2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}.struct_acc_t, %struct.{{.*}}.struct_acc_t* [[GEP_LAMBDA2]], i32 0, i32 0
7979
// CHECK: [[ARRAY_IDX2:%[a-zA-Z0-9_]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], [2 x [[ACCESSOR]]]* [[GEP_MEMBER_ACC2]], i64 0, i64 1
8080
// CHECK: [[MEM_LOAD2:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG1]].addr
8181
// CHECK: [[ACC_CAST2:%[0-9]+]] = addrspacecast [[ACCESSOR]]* [[ARRAY_IDX2]] to [[ACCESSOR]] addrspace(4)*
82-
// CHECK: call spir_func void @{{.*}}__init{{.*}}([[ACCESSOR]] addrspace(4)* [[ACC_CAST2]], i32 addrspace(1)* [[MEM_LOAD2]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE2]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE2]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET2]])
82+
// CHECK: call spir_func void @{{.*}}__init{{.*}}([[ACCESSOR]] addrspace(4)* {{[^,]*}} [[ACC_CAST2]], i32 addrspace(1)* [[MEM_LOAD2]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE2]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE2]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET2]])

‎clang/test/CodeGenSYCL/noexcept.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -44,7 +44,7 @@ void foo_cleanup() {
4444
// Regular + exception cleanup
4545
// CHECK-HOST-LIN: call void @_ZN1AD1Ev
4646
// CHECK-HOST-LIN: call void @_ZN1AD2Ev
47-
// CHECK-HOST-WIN: call void @"??1A@@QEAA@XZ"(%struct.A* %a)
47+
// CHECK-HOST-WIN: call void @"??1A@@QEAA@XZ"(%struct.A* {{[^,]*}} %a)
4848
}
4949

5050
template <typename name, typename Func>

‎clang/test/CodeGenSYCL/sampler.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -9,7 +9,7 @@
99
// CHECK-NEXT: [[GEP:%[0-9]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[ANON]], i32 0, i32 0
1010
// CHECK-NEXT: [[LOAD_SAMPLER_ARG:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)** [[SAMPLER_ARG]].addr, align 8
1111
// CHECK-NEXT: [[GEPCAST:%[0-9]+]] = addrspacecast %"class{{.*}}.cl::sycl::sampler"* [[GEP]] to %"class{{.*}}.cl::sycl::sampler" addrspace(4)*
12-
// CHECK-NEXT: call spir_func void @{{[a-zA-Z0-9_]+}}(%"class.{{.*}}.cl::sycl::sampler" addrspace(4)* [[GEPCAST]], %opencl.sampler_t addrspace(2)* [[LOAD_SAMPLER_ARG]])
12+
// CHECK-NEXT: call spir_func void @{{[a-zA-Z0-9_]+}}(%"class.{{.*}}.cl::sycl::sampler" addrspace(4)* {{[^,]*}} [[GEPCAST]], %opencl.sampler_t addrspace(2)* [[LOAD_SAMPLER_ARG]])
1313
//
1414

1515
// CHECK: define spir_kernel void @{{[a-zA-Z0-9_]+}}(%opencl.sampler_t addrspace(2)* [[SAMPLER_ARG_WRAPPED:%[a-zA-Z0-9_]+]], i32 [[ARG_A:%[a-zA-Z0-9_]+]])

‎clang/test/CodeGenSYCL/spir-calling-conv.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -9,9 +9,9 @@ int main() {
99

1010
// CHECK: define spir_kernel void @_ZTSZ4mainE15kernel_function()
1111

12-
// CHECK: call spir_func void @"_ZZ4mainENK3$_0clEv"(%"class.{{.*}}.anon" addrspace(4)* %2)
12+
// CHECK: call spir_func void @"_ZZ4mainENK3$_0clEv"(%"class.{{.*}}.anon" addrspace(4)* {{[^,]*}} %2)
1313

14-
// CHECK: define internal spir_func void @"_ZZ4mainENK3$_0clEv"(%"class.{{.*}}anon" addrspace(4)* %this)
14+
// CHECK: define internal spir_func void @"_ZZ4mainENK3$_0clEv"(%"class.{{.*}}anon" addrspace(4)* {{[^,]*}} %this)
1515

1616
kernel_single_task<class kernel_function>([]() {});
1717
return 0;

‎clang/test/CodeGenSYCL/spir-enum.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -23,7 +23,7 @@ int main() {
2323
// CHECK: define spir_kernel void @_ZTSZ4test9enum_typeE15kernel_function(i32 %_arg_)
2424

2525
// CHECK: getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"*
26-
// CHECK: call spir_func void @"_ZZ4test9enum_typeENK3$_0clEv"(%"class.{{.*}}.anon" addrspace(4)* %4)
26+
// CHECK: call spir_func void @"_ZZ4test9enum_typeENK3$_0clEv"(%"class.{{.*}}.anon" addrspace(4)* {{[^,]*}} %4)
2727

2828

2929
test( enum_type::B );

‎clang/test/CodeGenSYCL/stream.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -2,8 +2,8 @@
22
// RUN: FileCheck < %t.ll --enable-var-scope %s
33
//
44
// CHECK: define spir_kernel void @"{{.*}}StreamTester"(%"{{.*}}cl::sycl::stream"* byval(%"{{.*}}cl::sycl::stream") {{.*}}){{.*}}
5-
// CHECK: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::stream{{.*}} addrspace(4)* %{{[0-9]+}})
6-
// CHECK: call spir_func void @{{.*}}__finalize{{.*}}(%{{.*}}cl::sycl::stream{{.*}} addrspace(4)* %{{[0-9]+}})
5+
// CHECK: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::stream{{.*}} addrspace(4)* {{[^,]*}} %{{[0-9]+}})
6+
// CHECK: call spir_func void @{{.*}}__finalize{{.*}}(%{{.*}}cl::sycl::stream{{.*}} addrspace(4)* {{[^,]*}} %{{[0-9]+}})
77
//
88

99
#include "Inputs/sycl.hpp"

‎clang/test/CodeGenSYCL/sycl_kernel-host.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -11,10 +11,10 @@
1111
// CHECK: call spir_func void @{{.*}}KernelImpl{{.*}}({{.*}}, i32 1, double 2.000000e+00)
1212
// CHECK: define internal spir_func void @{{.*}}KernelImpl{{.*}}({{.*}} %f, i32 %i, double %d) #[[SKA]] {
1313
// CHECK: call spir_func void @"{{.*}}func{{.*}}"(%class
14-
// CHECK: define internal spir_func void @{{.*}}func{{.*}}(%class.anon* %this, i32 %i, double %d) #[[ALWAYSINLINE:[0-9]+]]
14+
// CHECK: define internal spir_func void @{{.*}}func{{.*}}(%class.anon* {{[^,]*}} %this, i32 %i, double %d) #[[ALWAYSINLINE:[0-9]+]]
1515
// CHECK: define linkonce_odr spir_func void @{{.*}}KernelImpl{{.*}}Functor{{.*}}({{.*}}, i32 %i, double %d) #[[SKA]] comdat {
1616
// CHECK: call spir_func void @{{.*}}Functor{{.*}}(%struct
17-
// CHECK: define linkonce_odr spir_func void @{{.*}}Functor{{.*}}(%struct.Functor* %this, i32 %i, double %d) #[[ALWAYSINLINE]]
17+
// CHECK: define linkonce_odr spir_func void @{{.*}}Functor{{.*}}(%struct.Functor* {{[^,]*}} %this, i32 %i, double %d) #[[ALWAYSINLINE]]
1818

1919
template <typename Func>
2020
void __attribute__((sycl_kernel))

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -38,4 +38,4 @@ int main() {
3838
// CHECK: [[MEMCPY_SRC:%[0-9a-zA-Z_]+]] = bitcast %union.{{.*}}MyUnion* [[MEM_ARG]] to i8*
3939
// CHECK: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 [[MEMCPY_DST]], i8* align 4 [[MEMCPY_SRC]], i64 12, i1 false)
4040
// CHECK: [[ACC_CAST1:%[0-9]+]] = addrspacecast %"class.{{.*}}.anon"* [[LOCAL_OBJECT]] to %"class.{{.*}}.anon" addrspace(4)*
41-
// CHECK: call spir_func void @{{.*}}(%"class.{{.*}}.anon" addrspace(4)* [[ACC_CAST1]])
41+
// CHECK: call spir_func void @{{.*}}(%"class.{{.*}}.anon" addrspace(4)* {{[^,]*}} [[ACC_CAST1]])

‎sycl/doc/GetStartedGuide.md

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -63,7 +63,7 @@ set DPCPP_HOME=%USERPROFILE%\sycl_workspace
6363
mkdir %DPCPP_HOME%
6464
cd %DPCPP_HOME%
6565
66-
git clone https://github.com/intel/llvm -b sycl
66+
git clone --config core.autocrlf=false https://github.com/intel/llvm -b sycl
6767
```
6868

6969
## Build DPC++ toolchain

‎sycl/include/CL/sycl/INTEL/esimd.hpp

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -10,6 +10,8 @@
1010

1111
#pragma once
1212

13+
/// \defgroup sycl_esimd DPC++ Explicit SIMD API
14+
1315
#include <CL/sycl/INTEL/esimd/esimd.hpp>
1416
#include <CL/sycl/INTEL/esimd/esimd_math.hpp>
1517
#include <CL/sycl/INTEL/esimd/esimd_memory.hpp>

0 commit comments

Comments
 (0)