Skip to content

Commit fa740a5

Browse files
author
Erich Keane
committed
Fix most of the CodeGenSYCL tests, int_header_sycl2020_spec_const.cpp crashes
1 parent a723f9a commit fa740a5

33 files changed

+133
-133
lines changed

‎clang/test/CodeGenSYCL/accessor_inheritance.cpp

+5-5
Original file line numberDiff line numberDiff line change
@@ -49,8 +49,8 @@ int main() {
4949
// CHECK: [[ARG_C]].addr.ascast = addrspacecast i32* [[ARG_C]].addr to i32 addrspace(4)*
5050
//
5151
// Lambda object alloca
52-
// CHECK: [[KERNEL:%[a-zA-Z0-9_]+]] = alloca %"class{{.*}}.anon"
53-
// CHECK: [[KERNEL_OBJ:%[a-zA-Z0-9_]+]] = addrspacecast %"class{{.*}}.anon"* [[KERNEL]] to %"class{{.*}}.anon" addrspace(4)*
52+
// CHECK: [[KERNEL:%[a-zA-Z0-9_]+]] = alloca %class{{.*}}.anon
53+
// CHECK: [[KERNEL_OBJ:%[a-zA-Z0-9_]+]] = addrspacecast %class{{.*}}.anon* [[KERNEL]] to %class{{.*}}.anon addrspace(4)*
5454
//
5555
// Kernel argument stores
5656
// CHECK: store i32 [[ARG_A]], i32 addrspace(4)* [[ARG_A]].addr.ascast
@@ -60,7 +60,7 @@ int main() {
6060
// CHECK: store i32 [[ARG_C]], i32 addrspace(4)* [[ARG_C]].addr.ascast
6161
//
6262
// Check A and B scalar fields initialization
63-
// CHECK: [[GEP:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class{{.*}}.anon", %"class{{.*}}.anon" addrspace(4)* [[KERNEL_OBJ]], i32 0, i32 0
63+
// CHECK: [[GEP:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.anon, %class{{.*}}.anon addrspace(4)* [[KERNEL_OBJ]], i32 0, i32 0
6464
// CHECK: [[BITCAST:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured addrspace(4)* [[GEP]] to %struct{{.*}}Base addrspace(4)*
6565
// CHECK: [[FIELD_A:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base addrspace(4)* [[BITCAST]], i32 0, i32 0
6666
// CHECK: [[ARG_A_LOAD:%[a-zA-Z0-9_]+]] = load i32, i32 addrspace(4)* [[ARG_A]].addr.ascast
@@ -85,13 +85,13 @@ int main() {
8585
// CHECK: store i32 [[ARG_C_LOAD]], i32 addrspace(4)* [[FIELD_C]]
8686
//
8787
// Check __init method calls
88-
// CHECK: [[GEP2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class{{.*}}.anon", %"class{{.*}}.anon" addrspace(4)* [[KERNEL_OBJ]], i32 0, i32 0
88+
// CHECK: [[GEP2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.anon, %class{{.*}}.anon addrspace(4)* [[KERNEL_OBJ]], i32 0, i32 0
8989
// CHECK: [[BITCAST3:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured addrspace(4)* [[GEP2]] to %struct{{.*}}Base addrspace(4)*
9090
// CHECK: [[ACC1_FIELD:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base addrspace(4)* [[BITCAST3]], i32 0, i32 2
9191
// CHECK: [[ACC1_DATA_LOAD:%[a-zA-Z0-9_]+]] = load i8 addrspace(1)*, i8 addrspace(1)* addrspace(4)* [[ACC1_DATA]].addr.ascast
9292
// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class{{.*}}cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[ACC1_FIELD]], i8 addrspace(1)* [[ACC1_DATA_LOAD]]
9393
//
94-
// CHECK: [[GEP3:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class{{.*}}.anon", %"class{{.*}}.anon" addrspace(4)* [[KERNEL_OBJ]], i32 0, i32 0
94+
// CHECK: [[GEP3:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class{{.*}}.anon, %class{{.*}}.anon addrspace(4)* [[KERNEL_OBJ]], i32 0, i32 0
9595
// CHECK: [[ACC2_DATA_LOAD:%[a-zA-Z0-9_]+]] = load i8 addrspace(1)*, i8 addrspace(1)* addrspace(4)* [[ACC2_DATA]].addr.ascast
9696
// CHECK: [[BITCAST4:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured addrspace(4)* [[GEP3]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)*
9797
// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class{{.*}}cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[BITCAST4]], i8 addrspace(1)* [[ACC2_DATA_LOAD]]

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

+4-4
Original file line numberDiff line numberDiff line change
@@ -27,8 +27,8 @@ int main() {
2727
// Check alloca for pointer argument
2828
// CHECK: [[MEM_ARG]].addr = alloca i32 addrspace(1)*
2929
// Check lambda object alloca
30-
// CHECK: [[ANONALLOCA:%[0-9]+]] = alloca %"class.{{.*}}.anon"
31-
// CHECK: [[ANON:%[0-9]+]] = addrspacecast %"class.{{.*}}.anon"* [[ANONALLOCA]] to %"class.{{.*}}.anon" addrspace(4)*
30+
// CHECK: [[ANONALLOCA:%[0-9]+]] = alloca %class.{{.*}}.anon
31+
// CHECK: [[ANON:%[0-9]+]] = addrspacecast %class.{{.*}}.anon* [[ANONALLOCA]] to %class.{{.*}}.anon addrspace(4)*
3232
// Check allocas for ranges
3333
// CHECK: [[ARANGEA:%agg.tmp.*]] = alloca %"struct.{{.*}}.cl::sycl::range"
3434
// CHECK: [[ARANGET:%agg.tmp.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::range"* [[ARANGEA]] to %"struct.{{.*}}.cl::sycl::range" addrspace(4)*
@@ -44,7 +44,7 @@ int main() {
4444
// CHECK: call spir_func {{.*}}accessor
4545

4646
// Check accessor GEP
47-
// CHECK: [[ACCESSOR:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon" addrspace(4)* [[ANON]], i32 0, i32 0
47+
// CHECK: [[ACCESSOR:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.{{.*}}.anon, %class.{{.*}}.anon addrspace(4)* [[ANON]], i32 0, i32 0
4848

4949
// Check load from kernel pointer argument alloca
5050
// CHECK: [[MEM_LOAD:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(4)* [[MEM_ARG]].addr.ascast
@@ -56,4 +56,4 @@ int main() {
5656
// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[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]])
5757

5858
// Check lambda "()" operator call
59-
// CHECK: call spir_func void @{{.*}}(%"class.{{.*}}.anon" addrspace(4)* {{[^,]*}})
59+
// CHECK: call spir_func void @{{.*}}(%class.{{.*}}.anon addrspace(4)* {{[^,]*}})

‎clang/test/CodeGenSYCL/debug-info-srcpos-kernel.cpp

+2-2
Original file line numberDiff line numberDiff line change
@@ -25,9 +25,9 @@ int main() {
2525
}
2626

2727
// CHECK: define{{.*}} spir_kernel {{.*}}19use_kernel_for_test({{.*}}){{.*}} !dbg [[KERNEL:![0-9]+]] {{.*}}{
28-
// CHECK: getelementptr inbounds %"class.{{.*}}.anon"{{.*}} !dbg [[LINE_A0:![0-9]+]]
28+
// CHECK: getelementptr inbounds %class.{{.*}}.anon{{.*}} !dbg [[LINE_A0:![0-9]+]]
2929
// CHECK: call spir_func void {{.*}}6__init{{.*}} !dbg [[LINE_A0]]
30-
// CHECK: call spir_func void @"_ZZ4mainENK3$_0clEv"{{.*}} !dbg [[LINE_B0:![0-9]+]]
30+
// CHECK: call spir_func void @_ZZ4mainENKUlvE_clEv{{.*}} !dbg [[LINE_B0:![0-9]+]]
3131
// CHECK: ret void, !dbg [[LINE_C0:![0-9]+]]
3232
// CHECK: [[FILE:![0-9]+]] = !DIFile(filename: "{{.*}}debug-info-srcpos-kernel.cpp"{{.*}})
3333
// CHECK: [[KERNEL]] = {{.*}}!DISubprogram(name: "{{.*}}19use_kernel_for_test"

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

+1-1
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 @_ZZ4mainENKUlvE_clEv(%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/device-variables.cpp

+2-2
Original file line numberDiff line numberDiff line change
@@ -32,10 +32,10 @@ int main() {
3232
// CHECK: store i32 1, i32 addrspace(4)* %b
3333
foo(local_value);
3434
// Local variables and constexprs captured by lambda
35-
// CHECK: [[GEP:%[0-9]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon" addrspace(4)* %{{.*}}, i32 0, i32 0
35+
// CHECK: [[GEP:%[0-9]+]] = getelementptr inbounds %class.{{.*}}.anon, %class.{{.*}}.anon addrspace(4)* %{{.*}}, i32 0, i32 0
3636
// CHECK: call spir_func void @{{.*}}foo{{.*}}(i32 addrspace(4)* align 4 dereferenceable(4) [[GEP]])
3737
int some_device_local_var = some_local_var;
38-
// CHECK: [[GEP1:%[0-9]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon" addrspace(4)* %{{.*}}, i32 0, i32 1
38+
// CHECK: [[GEP1:%[0-9]+]] = getelementptr inbounds %class.{{.*}}.anon, %class.{{.*}}.anon addrspace(4)* %{{.*}}, i32 0, i32 1
3939
// CHECK: [[LOAD1:%[0-9]+]] = load i32, i32 addrspace(4)* [[GEP1]]
4040
// CHECK: store i32 [[LOAD1]], i32 addrspace(4)* %some_device_local_var
4141
});

‎clang/test/CodeGenSYCL/disable_loop_pipelining.cpp

+3-3
Original file line numberDiff line numberDiff line change
@@ -29,8 +29,8 @@ int main() {
2929
return 0;
3030
}
3131

32-
// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel1"() #0 {{.*}} !disable_loop_pipelining ![[NUM5:[0-9]+]]
33-
// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel2"() #0 {{.*}} ![[NUM4:[0-9]+]]
34-
// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel3"() #0 {{.*}} !disable_loop_pipelining ![[NUM5]]
32+
// CHECK: define dso_local spir_kernel void @{{.*}}test_kernel1() #0 {{.*}} !disable_loop_pipelining ![[NUM5:[0-9]+]]
33+
// CHECK: define dso_local spir_kernel void @{{.*}}test_kernel2() #0 {{.*}} ![[NUM4:[0-9]+]]
34+
// CHECK: define dso_local spir_kernel void @{{.*}}test_kernel3() #0 {{.*}} !disable_loop_pipelining ![[NUM5]]
3535
// CHECK: ![[NUM4]] = !{}
3636
// CHECK: ![[NUM5]] = !{i32 1}

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

+1-1
Original file line numberDiff line numberDiff line change
@@ -26,7 +26,7 @@ void test(int val) {
2626
});
2727

2828
// --- Name
29-
// CHECK-LABEL: define {{.*}}spir_kernel void @"_ZTSZZ4testiENK3$_0clERN2cl4sycl7handlerEE12esimd_kernel"(
29+
// CHECK-LABEL: define {{.*}}spir_kernel void @_ZTSZZ4testiENKUlRN2cl4sycl7handlerEE_clES2_E12esimd_kernel(
3030
// --- Attributes
3131
// CHECK: {{.*}} !kernel_arg_accessor_ptr ![[ACC_PTR_ATTR:[0-9]+]] !sycl_explicit_simd !{{[0-9]+}} {{.*}}{
3232
// --- init_esimd call is expected instead of __init:

‎clang/test/CodeGenSYCL/initiation_interval.cpp

+4-4
Original file line numberDiff line numberDiff line change
@@ -39,10 +39,10 @@ int main() {
3939
return 0;
4040
}
4141

42-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1"() #0 {{.*}} !initiation_interval ![[NUM1:[0-9]+]]
43-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2"() #0 {{.*}} !initiation_interval ![[NUM42:[0-9]+]]
44-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3"() #0 {{.*}} !initiation_interval ![[NUM2:[0-9]+]]
45-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4"() #0 {{.*}} ![[NUM0:[0-9]+]]
42+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() #0 {{.*}} !initiation_interval ![[NUM1:[0-9]+]]
43+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2() #0 {{.*}} !initiation_interval ![[NUM42:[0-9]+]]
44+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3() #0 {{.*}} !initiation_interval ![[NUM2:[0-9]+]]
45+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() #0 {{.*}} ![[NUM0:[0-9]+]]
4646
// CHECK: ![[NUM0]] = !{}
4747
// CHECK: ![[NUM1]] = !{i32 1}
4848
// CHECK: ![[NUM42]] = !{i32 42}

‎clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp

+5-5
Original file line numberDiff line numberDiff line change
@@ -40,10 +40,10 @@ int main() {
4040
return 0;
4141
}
4242

43-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1"() #0 {{.*}} !no_global_work_offset ![[NUM5:[0-9]+]]
44-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2"() #0 {{.*}} !no_global_work_offset ![[NUM5]]
45-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3"() #0 {{.*}} ![[NUM4:[0-9]+]]
46-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4"() #0 {{.*}} !no_global_work_offset ![[NUM5]]
47-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5"() #0 {{.*}} !no_global_work_offset ![[NUM5]]
43+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() #0 {{.*}} !no_global_work_offset ![[NUM5:[0-9]+]]
44+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2() #0 {{.*}} !no_global_work_offset ![[NUM5]]
45+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3() #0 {{.*}} ![[NUM4:[0-9]+]]
46+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() #0 {{.*}} !no_global_work_offset ![[NUM5]]
47+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5() #0 {{.*}} !no_global_work_offset ![[NUM5]]
4848
// CHECK-NOT: ![[NUM4]] = !{i32 0}
4949
// CHECK: ![[NUM5]] = !{}

‎clang/test/CodeGenSYCL/intel-max-global-work-dim.cpp

+4-4
Original file line numberDiff line numberDiff line change
@@ -37,9 +37,9 @@ int main() {
3737
return 0;
3838
}
3939

40-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1"() #0 {{.*}} !max_global_work_dim ![[NUM1:[0-9]+]]
41-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2"() #0 {{.*}} !max_global_work_dim ![[NUM2:[0-9]+]]
42-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3"() #0 {{.*}} !max_global_work_dim ![[NUM2]]
43-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4"() #0 {{.*}} !max_global_work_dim ![[NUM2]]
40+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() #0 {{.*}} !max_global_work_dim ![[NUM1:[0-9]+]]
41+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2() #0 {{.*}} !max_global_work_dim ![[NUM2:[0-9]+]]
42+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3() #0 {{.*}} !max_global_work_dim ![[NUM2]]
43+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() #0 {{.*}} !max_global_work_dim ![[NUM2]]
4444
// CHECK: ![[NUM1]] = !{i32 1}
4545
// CHECK: ![[NUM2]] = !{i32 2}

‎clang/test/CodeGenSYCL/intel-max-work-group-size.cpp

+5-5
Original file line numberDiff line numberDiff line change
@@ -45,11 +45,11 @@ int main() {
4545
return 0;
4646
}
4747

48-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1"() #0 {{.*}} !max_work_group_size ![[NUM1:[0-9]+]]
49-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2"() #0 {{.*}} !max_work_group_size ![[NUM8:[0-9]+]]
50-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3"() #0 {{.*}} !max_work_group_size ![[NUM6:[0-9]+]]
51-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4"() #0 {{.*}} !max_work_group_size ![[NUM2:[0-9]+]]
52-
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5"() #0 {{.*}} !max_work_group_size ![[NUM4:[0-9]+]]
48+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() #0 {{.*}} !max_work_group_size ![[NUM1:[0-9]+]]
49+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2() #0 {{.*}} !max_work_group_size ![[NUM8:[0-9]+]]
50+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3() #0 {{.*}} !max_work_group_size ![[NUM6:[0-9]+]]
51+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() #0 {{.*}} !max_work_group_size ![[NUM2:[0-9]+]]
52+
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5() #0 {{.*}} !max_work_group_size ![[NUM4:[0-9]+]]
5353
// CHECK: ![[NUM1]] = !{i32 1, i32 1, i32 1}
5454
// CHECK: ![[NUM8]] = !{i32 8, i32 8, i32 8}
5555
// CHECK: ![[NUM6]] = !{i32 6, i32 3, i32 1}

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

+3-3
Original file line numberDiff line numberDiff line change
@@ -22,16 +22,16 @@ void test(int val) {
2222
});
2323
}
2424

25-
// NONATIVESUPPORT: define dso_local void @"{{.*}}test_kernel_handler{{.*}}"
25+
// NONATIVESUPPORT: define dso_local void @{{.*}}test_kernel_handler{{[^(]*}}
2626
// NONATIVESUPPORT-SAME: (i32 %_arg_, i8 addrspace(1)* %_arg__specialization_constants_buffer)
2727
// NONATIVESUPPORT: %kh = alloca %"class.[[MANGLEDCLASS:[a-zA-Z0-9_]+]].cl::sycl::kernel_handler", align 1
2828
// NONATIVESUPPORT: %[[KH:[0-9]+]] = load i8 addrspace(1)*, i8 addrspace(1)** %_arg__specialization_constants_buffer.addr, align 8
2929
// NONATIVESUPPORT: %[[ADDRSPACECAST:[0-9]+]] = addrspacecast i8 addrspace(1)* %[[KH]] to i8*
3030
// NONATIVESUPPORT: call void @{{.*}}__init_specialization_constants_buffer{{.*}}(%"class.[[MANGLEDCLASS]].cl::sycl::kernel_handler"* nonnull align 1 dereferenceable(1) %kh, i8* %[[ADDRSPACECAST]])
31-
// NONATIVESUPPORT: void @"[[MANGLEDKERNELCALL:[a-zA-Z0-9_$]+]]"
31+
// NONATIVESUPPORT: void @[[MANGLEDKERNELCALL:[a-zA-Z0-9_$]+]]
3232
// NONATIVESUPPORT-SAME: byval(%"class.[[MANGLEDCLASS]].cl::sycl::kernel_handler")
3333

34-
// NATIVESUPPORT: define dso_local spir_kernel void @"{{.*}}test_kernel_handler{{.*}}"
34+
// NATIVESUPPORT: define dso_local spir_kernel void @{{.*}}test_kernel_handler{{[^(]*}}
3535
// NATIVESUPPORT-SAME: (i32 %_arg_)
3636
// NATIVESUPPORT: %kh = alloca %"class.[[MANGLEDCLASS:[a-zA-Z0-9_]+]].cl::sycl::kernel_handler"
3737
// NATIVESUPPORT-NOT: __init_specialization_constants_buffer

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

+5-5
Original file line numberDiff line numberDiff line change
@@ -39,8 +39,8 @@ int main() {
3939
// CHECK: [[MEM_ARG2:%[a-zA-Z0-9_.]+]] = alloca i32 addrspace(1)*, align 8
4040

4141
// CHECK lambda object alloca
42-
// CHECK: [[LOCAL_OBJECTA:%0]] = alloca %"class.{{.*}}.anon", align 4
43-
// CHECK: [[LOCAL_OBJECT:%.*]] = addrspacecast %"class.{{.*}}.anon"* [[LOCAL_OBJECTA]] to %"class.{{.*}}.anon" addrspace(4)*
42+
// CHECK: [[LOCAL_OBJECTA:%0]] = alloca %class.{{.*}}.anon, align 4
43+
// CHECK: [[LOCAL_OBJECT:%.*]] = addrspacecast %class.{{.*}}.anon* [[LOCAL_OBJECTA]] to %class.{{.*}}.anon addrspace(4)*
4444

4545
// CHECK allocas for ranges
4646
// CHECK: [[ACC_RANGE1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range"
@@ -57,7 +57,7 @@ int main() {
5757
// CHECK: [[OFFSET2AS:%.*]] = addrspacecast %"struct.{{.*}}.cl::sycl::id"* [[OFFSET2A]] to %"struct.{{.*}}.cl::sycl::id" addrspace(4)*
5858

5959
// CHECK accessor array default inits
60-
// CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon" addrspace(4)* [[LOCAL_OBJECT]], i32 0, i32 0
60+
// CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.{{.*}}.anon, %class.{{.*}}.anon addrspace(4)* [[LOCAL_OBJECT]], i32 0, i32 0
6161
// CHECK: [[BEGIN:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR:.*]]], [2 x [[ACCESSOR]]] addrspace(4)* [[ACCESSOR_ARRAY1]], i64 0, i64 0
6262
// Clang takes advantage of element 1 having the same address as the array, so it doesn't do a GEP.
6363
// CTOR Call #1
@@ -67,7 +67,7 @@ int main() {
6767
// CHECK: call spir_func void @{{.+}}([[ACCESSOR]] addrspace(4)* {{[^,]*}} [[ELEM2_GEP]])
6868

6969
// CHECK acc[0] __init method call
70-
// CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon" addrspace(4)* [[LOCAL_OBJECT]], i32 0, i32 0
70+
// CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.{{.*}}.anon, %class.{{.*}}.anon addrspace(4)* [[LOCAL_OBJECT]], i32 0, i32 0
7171
// CHECK: [[INDEX1:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], [2 x [[ACCESSOR]]] addrspace(4)* [[ACCESSOR_ARRAY1]], i64 0, i64 0
7272
// CHECK load from kernel pointer argument alloca
7373
// CHECK: [[MEM_LOAD1:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(4)* [[MEM_ARG1]]
@@ -77,7 +77,7 @@ int main() {
7777
// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[INDEX1]], 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]])
7878

7979
// CHECK acc[1] __init method call
80-
// CHECK: [[ACCESSOR_ARRAY2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon" addrspace(4)* [[LOCAL_OBJECT]], i32 0, i32 0
80+
// CHECK: [[ACCESSOR_ARRAY2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.{{.*}}.anon, %class.{{.*}}.anon addrspace(4)* [[LOCAL_OBJECT]], i32 0, i32 0
8181
// CHECK: [[INDEX2:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], [2 x [[ACCESSOR]]] addrspace(4)* [[ACCESSOR_ARRAY2]], i64 0, i64 1
8282
// CHECK load from kernel pointer argument alloca
8383
// CHECK: [[MEM_LOAD2:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(4)* [[MEM_ARG2]]

0 commit comments

Comments
 (0)