1
1
// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
2
+ // Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --include-generated-funcs --version 5
3
+ // And edited to fix some failures.
2
4
3
5
// This test checks a kernel argument that is an Accessor array
4
6
@@ -22,69 +24,58 @@ int main() {
22
24
acc[1 ].use ();
23
25
});
24
26
}
25
-
26
- // Check kernel_A parameters
27
- // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_A
28
- // CHECK-SAME: ptr addrspace(1) noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]],
29
- // CHECK-SAME: ptr noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]],
30
- // CHECK-SAME: ptr noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]],
31
- // CHECK-SAME: ptr noundef byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]],
32
- // CHECK-SAME: ptr addrspace(1) noundef align 4 [[MEM_ARG2:%[a-zA-Z0-9_]+4]],
33
- // CHECK-SAME: ptr noundef byval{{.*}}align 4 [[ACC_RANGE2:%[a-zA-Z0-9_]+6]],
34
- // CHECK-SAME: ptr noundef byval{{.*}}align 4 [[MEM_RANGE2:%[a-zA-Z0-9_]+7]],
35
- // CHECK-SAME: ptr noundef byval{{.*}}align 4 [[OFFSET2:%[a-zA-Z0-9_]+8]])
36
-
37
- // CHECK alloca for pointer arguments
38
- // CHECK: [[MEM_ARG1:%[a-zA-Z0-9_.]+]] = alloca ptr addrspace(1), align 8
39
- // CHECK: [[MEM_ARG2:%[a-zA-Z0-9_.]+]] = alloca ptr addrspace(1), align 8
40
-
41
- // CHECK lambda object alloca
42
- // CHECK: [[LOCAL_OBJECTA:%__SYCLKernel]] = alloca %class.anon, align 4
43
-
44
- // CHECK allocas for ranges
45
- // CHECK: [[ACC_RANGE1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.sycl::_V1::range"
46
- // CHECK: [[MEM_RANGE1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.sycl::_V1::range"
47
- // CHECK: [[OFFSET1A:%[a-zA-Z0-9_.]+]] = alloca %"struct.sycl::_V1::id"
48
- // CHECK: [[ACC_RANGE2A:%[a-zA-Z0-9_.]+]] = alloca %"struct.sycl::_V1::range"
49
- // CHECK: [[MEM_RANGE2A:%[a-zA-Z0-9_.]+]] = alloca %"struct.sycl::_V1::range"
50
- // CHECK: [[OFFSET2A:%[a-zA-Z0-9_.]+]] = alloca %"struct.sycl::_V1::id"
51
-
52
- // CHECK lambda object addrspacecast
53
- // CHECK: [[LOCAL_OBJECT:%.*]] = addrspacecast ptr [[LOCAL_OBJECTA]] to ptr addrspace(4)
54
-
55
- // CHECK addrspacecasts for ranges
56
- // CHECK: [[ACC_RANGE1AS:%.*]] = addrspacecast ptr [[ACC_RANGE1A]] to ptr addrspace(4)
57
- // CHECK: [[MEM_RANGE1AS:%.*]] = addrspacecast ptr [[MEM_RANGE1A]] to ptr addrspace(4)
58
- // CHECK: [[OFFSET1AS:%.*]] = addrspacecast ptr [[OFFSET1A]] to ptr addrspace(4)
59
- // CHECK: [[ACC_RANGE2AS:%.*]] = addrspacecast ptr [[ACC_RANGE2A]] to ptr addrspace(4)
60
- // CHECK: [[MEM_RANGE2AS:%.*]] = addrspacecast ptr [[MEM_RANGE2A]] to ptr addrspace(4)
61
- // CHECK: [[OFFSET2AS:%.*]] = addrspacecast ptr [[OFFSET2A]] to ptr addrspace(4)
62
- // CHECK accessor array default inits
63
- // CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.anon, ptr addrspace(4) [[LOCAL_OBJECT]], i32 0, i32 0
64
- // CHECK: [[BEGIN:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR:.*]]], ptr addrspace(4) [[ACCESSOR_ARRAY1]], i64 0, i64 0
65
- // Clang takes advantage of element 1 having the same address as the array, so it doesn't do a GEP.
66
- // CTOR Call #1
67
- // CHECK: call spir_func void @{{.+}}(ptr addrspace(4) {{[^,]*}} [[BEGIN]])
68
- // CHECK: [[ELEM2_GEP:%[a-zA-Z0-9_.]+]] = getelementptr inbounds [[ACCESSOR]], ptr addrspace(4) [[BEGIN]], i64 1
69
- // CTOR Call #2
70
- // CHECK: call spir_func void @{{.+}}(ptr addrspace(4) {{[^,]*}} [[ELEM2_GEP]])
71
-
72
- // CHECK acc[0] __init method call
73
- // CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.anon, ptr addrspace(4) [[LOCAL_OBJECT]], i32 0, i32 0
74
- // CHECK: [[INDEX1:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], ptr addrspace(4) [[ACCESSOR_ARRAY1]], i64 0, i64 0
75
- // CHECK load from kernel pointer argument alloca
76
- // CHECK: [[MEM_LOAD1:%[a-zA-Z0-9_]+]] = load ptr addrspace(1), ptr addrspace(4) [[MEM_ARG1]]
77
- // CHECK: [[ACC_RANGE1:%.*]] = addrspacecast ptr addrspace(4) [[ACC_RANGE1AS]] to ptr
78
- // CHECK: [[MEM_RANGE1:%.*]] = addrspacecast ptr addrspace(4) [[MEM_RANGE1AS]] to ptr
79
- // CHECK: [[OFFSET1:%.*]] = addrspacecast ptr addrspace(4) [[OFFSET1AS]] to ptr
80
- // CHECK: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) {{[^,]*}} [[INDEX1]], ptr addrspace(1) noundef [[MEM_LOAD1]], ptr noundef byval({{.*}}) align 4 [[ACC_RANGE1]], ptr noundef byval({{.*}}) align 4 [[MEM_RANGE1]], ptr noundef byval({{.*}}) align 4 [[OFFSET1]])
81
-
82
- // CHECK acc[1] __init method call
83
- // CHECK: [[ACCESSOR_ARRAY2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.anon, ptr addrspace(4) [[LOCAL_OBJECT]], i32 0, i32 0
84
- // CHECK: [[INDEX2:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], ptr addrspace(4) [[ACCESSOR_ARRAY2]], i64 0, i64 1
85
- // CHECK load from kernel pointer argument alloca
86
- // CHECK: [[MEM_LOAD2:%[a-zA-Z0-9_]+]] = load ptr addrspace(1), ptr addrspace(4) [[MEM_ARG2]]
87
- // CHECK: [[ACC_RANGE2:%.*]] = addrspacecast ptr addrspace(4) [[ACC_RANGE2AS]] to ptr
88
- // CHECK: [[MEM_RANGE2:%.*]] = addrspacecast ptr addrspace(4) [[MEM_RANGE2AS]] to ptr
89
- // CHECK: [[OFFSET2:%.*]] = addrspacecast ptr addrspace(4) [[OFFSET2AS]] to ptr
90
- // CHECK: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) {{[^,]*}} [[INDEX2]], ptr addrspace(1) noundef [[MEM_LOAD2]], ptr noundef byval({{.*}}) align 4 [[ACC_RANGE2]], ptr noundef byval({{.*}}) align 4 [[MEM_RANGE2]], ptr noundef byval({{.*}}) align 4 [[OFFSET2]])
27
+ // CHECK-LABEL: define dso_local spir_kernel void @_ZTSZ4mainE8kernel_A(
28
+ // 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]+]]
29
+ // CHECK-NEXT: [[ENTRY:.*:]]
30
+ // CHECK-NEXT: [[_ARG_ACC_ADDR:%.*]] = alloca ptr addrspace(1), align 8
31
+ // CHECK-NEXT: [[_ARG_ACC_ADDR5:%.*]] = alloca ptr addrspace(1), align 8
32
+ // CHECK-NEXT: [[__SYCLKERNEL:%.*]] = alloca [[CLASS_ANON:%.*]], align 4
33
+ // CHECK-NEXT: [[AGG_TMP:%.*]] = alloca %"struct.sycl::_V1::range", align 4
34
+ // CHECK-NEXT: [[AGG_TMP10:%.*]] = alloca %"struct.sycl::_V1::range", align 4
35
+ // CHECK-NEXT: [[AGG_TMP11:%.*]] = alloca %"struct.sycl::_V1::id", align 4
36
+ // CHECK-NEXT: [[AGG_TMP14:%.*]] = alloca %"struct.sycl::_V1::range", align 4
37
+ // CHECK-NEXT: [[AGG_TMP15:%.*]] = alloca %"struct.sycl::_V1::range", align 4
38
+ // CHECK-NEXT: [[AGG_TMP16:%.*]] = alloca %"struct.sycl::_V1::id", align 4
39
+ // CHECK-NEXT: [[_ARG_ACC_ADDR_ASCAST:%.*]] = addrspacecast ptr [[_ARG_ACC_ADDR]] to ptr addrspace(4)
40
+ // CHECK-NEXT: [[_ARG_ACC_ADDR5_ASCAST:%.*]] = addrspacecast ptr [[_ARG_ACC_ADDR5]] to ptr addrspace(4)
41
+ // CHECK-NEXT: [[__SYCLKERNEL_ASCAST:%.*]] = addrspacecast ptr [[__SYCLKERNEL]] to ptr addrspace(4)
42
+ // CHECK-NEXT: [[AGG_TMP_ASCAST:%.*]] = addrspacecast ptr [[AGG_TMP]] to ptr addrspace(4)
43
+ // CHECK-NEXT: [[AGG_TMP10_ASCAST:%.*]] = addrspacecast ptr [[AGG_TMP10]] to ptr addrspace(4)
44
+ // CHECK-NEXT: [[AGG_TMP11_ASCAST:%.*]] = addrspacecast ptr [[AGG_TMP11]] to ptr addrspace(4)
45
+ // CHECK-NEXT: [[AGG_TMP14_ASCAST:%.*]] = addrspacecast ptr [[AGG_TMP14]] to ptr addrspace(4)
46
+ // CHECK-NEXT: [[AGG_TMP15_ASCAST:%.*]] = addrspacecast ptr [[AGG_TMP15]] to ptr addrspace(4)
47
+ // CHECK-NEXT: [[AGG_TMP16_ASCAST:%.*]] = addrspacecast ptr [[AGG_TMP16]] to ptr addrspace(4)
48
+ // CHECK-NEXT: store ptr addrspace(1) [[_ARG_ACC]], ptr addrspace(4) [[_ARG_ACC_ADDR_ASCAST]], align 8
49
+ // CHECK-NEXT: [[_ARG_ACC1_ASCAST:%.*]] = addrspacecast ptr [[_ARG_ACC1]] to ptr addrspace(4)
50
+ // CHECK-NEXT: [[_ARG_ACC2_ASCAST:%.*]] = addrspacecast ptr [[_ARG_ACC2]] to ptr addrspace(4)
51
+ // CHECK-NEXT: [[_ARG_ACC3_ASCAST:%.*]] = addrspacecast ptr [[_ARG_ACC3]] to ptr addrspace(4)
52
+ // CHECK-NEXT: store ptr addrspace(1) [[_ARG_ACC4]], ptr addrspace(4) [[_ARG_ACC_ADDR5_ASCAST]], align 8
53
+ // CHECK-NEXT: [[_ARG_ACC6_ASCAST:%.*]] = addrspacecast ptr [[_ARG_ACC6]] to ptr addrspace(4)
54
+ // CHECK-NEXT: [[_ARG_ACC7_ASCAST:%.*]] = addrspacecast ptr [[_ARG_ACC7]] to ptr addrspace(4)
55
+ // CHECK-NEXT: [[_ARG_ACC8_ASCAST:%.*]] = addrspacecast ptr [[_ARG_ACC8]] to ptr addrspace(4)
56
+ // CHECK-NEXT: [[ACC:%.*]] = getelementptr inbounds [[CLASS_ANON]], ptr addrspace(4) [[__SYCLKERNEL_ASCAST]], i32 0, i32 0
57
+ // CHECK-NEXT: call spir_func void @_ZN4sycl3_V18accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEC1Ev(ptr addrspace(4) noundef align 4 dereferenceable_or_null(12) [[ACC]]) #[[ATTR4:[0-9]+]]
58
+ // CHECK-NEXT: [[ARRAYINIT_ELEMENT:%.*]] = getelementptr inbounds %"class.sycl::_V1::accessor", ptr addrspace(4) [[ACC]], i64 1
59
+ // CHECK-NEXT: call spir_func void @_ZN4sycl3_V18accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEC1Ev(ptr addrspace(4) noundef align 4 dereferenceable_or_null(12) [[ARRAYINIT_ELEMENT]]) #[[ATTR4]]
60
+ // CHECK-NEXT: [[ACC9:%.*]] = getelementptr inbounds [[CLASS_ANON]], ptr addrspace(4) [[__SYCLKERNEL_ASCAST]], i32 0, i32 0
61
+ // CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [2 x %"class.sycl::_V1::accessor"], ptr addrspace(4) [[ACC9]], i64 0, i64 0
62
+ // CHECK-NEXT: [[TMP0:%.*]] = load ptr addrspace(1), ptr addrspace(4) [[_ARG_ACC_ADDR_ASCAST]], align 8
63
+ // CHECK-NEXT: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 4 [[AGG_TMP_ASCAST]], ptr addrspace(4) align 4 [[_ARG_ACC1_ASCAST]], i64 4, i1 false)
64
+ // CHECK-NEXT: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 4 [[AGG_TMP10_ASCAST]], ptr addrspace(4) align 4 [[_ARG_ACC2_ASCAST]], i64 4, i1 false)
65
+ // CHECK-NEXT: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 4 [[AGG_TMP11_ASCAST]], ptr addrspace(4) align 4 [[_ARG_ACC3_ASCAST]], i64 4, i1 false)
66
+ // CHECK-NEXT: [[AGG_TMP_ASCAST_ASCAST:%.*]] = addrspacecast ptr addrspace(4) [[AGG_TMP_ASCAST]] to ptr
67
+ // CHECK-NEXT: [[AGG_TMP10_ASCAST_ASCAST:%.*]] = addrspacecast ptr addrspace(4) [[AGG_TMP10_ASCAST]] to ptr
68
+ // CHECK-NEXT: [[AGG_TMP11_ASCAST_ASCAST:%.*]] = addrspacecast ptr addrspace(4) [[AGG_TMP11_ASCAST]] to ptr
69
+ // CHECK-NEXT: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) noundef align 4 dereferenceable_or_null(12) [[ARRAYIDX]], ptr addrspace(1) noundef [[TMP0]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[AGG_TMP_ASCAST_ASCAST]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[AGG_TMP10_ASCAST_ASCAST]], ptr noundef byval(%"struct.sycl::_V1::id") align 4 [[AGG_TMP11_ASCAST_ASCAST]]) #[[ATTR4]]
70
+ // CHECK-NEXT: [[ACC12:%.*]] = getelementptr inbounds [[CLASS_ANON]], ptr addrspace(4) [[__SYCLKERNEL_ASCAST]], i32 0, i32 0
71
+ // CHECK-NEXT: [[ARRAYIDX13:%.*]] = getelementptr inbounds [2 x %"class.sycl::_V1::accessor"], ptr addrspace(4) [[ACC12]], i64 0, i64 1
72
+ // CHECK-NEXT: [[TMP1:%.*]] = load ptr addrspace(1), ptr addrspace(4) [[_ARG_ACC_ADDR5_ASCAST]], align 8
73
+ // CHECK-NEXT: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 4 [[AGG_TMP14_ASCAST]], ptr addrspace(4) align 4 [[_ARG_ACC6_ASCAST]], i64 4, i1 false)
74
+ // CHECK-NEXT: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 4 [[AGG_TMP15_ASCAST]], ptr addrspace(4) align 4 [[_ARG_ACC7_ASCAST]], i64 4, i1 false)
75
+ // CHECK-NEXT: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 4 [[AGG_TMP16_ASCAST]], ptr addrspace(4) align 4 [[_ARG_ACC8_ASCAST]], i64 4, i1 false)
76
+ // CHECK-NEXT: [[AGG_TMP14_ASCAST_ASCAST:%.*]] = addrspacecast ptr addrspace(4) [[AGG_TMP14_ASCAST]] to ptr
77
+ // CHECK-NEXT: [[AGG_TMP15_ASCAST_ASCAST:%.*]] = addrspacecast ptr addrspace(4) [[AGG_TMP15_ASCAST]] to ptr
78
+ // CHECK-NEXT: [[AGG_TMP16_ASCAST_ASCAST:%.*]] = addrspacecast ptr addrspace(4) [[AGG_TMP16_ASCAST]] to ptr
79
+ // CHECK-NEXT: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) noundef align 4 dereferenceable_or_null(12) [[ARRAYIDX13]], ptr addrspace(1) noundef [[TMP1]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[AGG_TMP14_ASCAST_ASCAST]], ptr noundef byval(%"struct.sycl::_V1::range") align 4 [[AGG_TMP15_ASCAST_ASCAST]], ptr noundef byval(%"struct.sycl::_V1::id") align 4 [[AGG_TMP16_ASCAST_ASCAST]]) #[[ATTR4]]
80
+ // CHECK-NEXT: call spir_func void @_ZZ4mainENKUlvE_clEv(ptr addrspace(4) noundef align 4 dereferenceable_or_null(24) [[__SYCLKERNEL_ASCAST]]) #[[ATTR4]]
81
+ // CHECK-NEXT: ret void
0 commit comments