Skip to content

Commit 12fb6e9

Browse files
jsjijsji
authored
[Opaque pointer] Port sycl-post-link and SPIRITTAnnotations tests to opaque pointers (#10893)
- Port sycl-post-link to opaque pointers - Update SPIRITTAnnotations tests to opaque pointers
1 parent 7aed0ba commit 12fb6e9

File tree

15 files changed

+270
-283
lines changed

15 files changed

+270
-283
lines changed

‎llvm/test/Transforms/SPIRITTAnnotations/itt_atomic_load.ll

Lines changed: 53 additions & 53 deletions
Original file line numberDiff line numberDiff line change
@@ -24,38 +24,38 @@ $_ZTS11load_kernelIiE = comdat any
2424
@__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
2525

2626
; Function Attrs: convergent norecurse
27-
define weak_odr dso_local spir_kernel void @_ZTSN2cl4sycl6detail19__pf_kernel_wrapperI11load_kernelIiEEE(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_, i32 addrspace(1)* %_arg_1, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_3, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_4, i32 addrspace(1)* %_arg_5, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_7, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_8, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_9) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !4 {
27+
define weak_odr dso_local spir_kernel void @_ZTSN2cl4sycl6detail19__pf_kernel_wrapperI11load_kernelIiEEE(ptr byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_, ptr addrspace(1) %_arg_1, ptr byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, ptr byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_3, ptr byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_4, ptr addrspace(1) %_arg_5, ptr byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_7, ptr byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_8, ptr byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_9) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !4 {
2828
entry:
2929
; CHECK-LABEL: _ZTSN2cl4sycl6detail19__pf_kernel_wrapperI11load_kernelIiEEE(
3030
; CHECK-NEXT: entry:
3131
; CHECK-NEXT: call void @__itt_offload_wi_start_wrapper()
32-
%0 = getelementptr inbounds %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range", %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* %_arg_, i64 0, i32 0, i32 0, i64 0
33-
%1 = addrspacecast i64* %0 to i64 addrspace(4)*
34-
%2 = load i64, i64 addrspace(4)* %1, align 8
35-
%3 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId to <3 x i64> addrspace(4)*), align 32, !noalias !5
32+
%0 = getelementptr inbounds %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range", ptr %_arg_, i64 0, i32 0, i32 0, i64 0
33+
%1 = addrspacecast i64* %0 to ptr addrspace(4)
34+
%2 = load i64, ptr addrspace(4) %1, align 8
35+
%3 = load <3 x i64>, ptr addrspace(4) addrspacecast (ptr addrspace(1) @__spirv_BuiltInGlobalInvocationId to ptr addrspace(4)), align 32, !noalias !5
3636
%4 = extractelement <3 x i64> %3, i64 0
3737
%cmp.not.i = icmp ult i64 %4, %2
3838
br i1 %cmp.not.i, label %if.end.i, label %_ZZN2cl4sycl7handler24parallel_for_lambda_implI11load_kernelIiEZZ9load_testIiEvNS0_5queueEmENKUlRS1_E_clES7_EUlNS0_4itemILi1ELb1EEEE_Li1EEEvNS0_5rangeIXT1_EEET0_ENKUlSA_E_clESA_.exit
3939

4040
if.end.i: ; preds = %entry
41-
%5 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_9, i64 0, i32 0, i32 0, i64 0
42-
%6 = addrspacecast i64* %5 to i64 addrspace(4)*
43-
%7 = load i64, i64 addrspace(4)* %6, align 8
44-
%add.ptr.i = getelementptr inbounds i32, i32 addrspace(1)* %_arg_5, i64 %7
45-
%8 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_4, i64 0, i32 0, i32 0, i64 0
46-
%9 = addrspacecast i64* %8 to i64 addrspace(4)*
47-
%10 = load i64, i64 addrspace(4)* %9, align 8
48-
%add.ptr.i34 = getelementptr inbounds i32, i32 addrspace(1)* %_arg_1, i64 %10
49-
; CHECK: [[ARG_ASCAST:%[0-9a-zA-Z._]+]] = addrspacecast i32 addrspace(1)* %[[ATOMIC_ARG_1:[0-9a-zA-Z._]+]] to i8 addrspace(4)*
50-
; CHECK-NEXT: call void @__itt_offload_atomic_op_start(i8 addrspace(4)* [[ARG_ASCAST]], i32 0, i32 0)
51-
; CHECK-NEXT: {{.*}}__spirv_AtomicLoad{{.*}}(i32 addrspace(1)* %[[ATOMIC_ARG_1]],{{.*}}, i32 896
52-
; CHECK-NEXT: [[ARG_ASCAST:%[0-9a-zA-Z._]+]] = addrspacecast i32 addrspace(1)* %[[ATOMIC_ARG_1]] to i8 addrspace(4)*
53-
; CHECK-NEXT: call void @__itt_offload_atomic_op_finish(i8 addrspace(4)* [[ARG_ASCAST]], i32 0, i32 0)
54-
%call3.i.i.i.i = tail call spir_func i32 @_Z18__spirv_AtomicLoadPU3AS1KiN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE(i32 addrspace(1)* %add.ptr.i34, i32 1, i32 896) #2
55-
call spir_func void @__synthetic_spir_fun_call(i32 addrspace(1)* %add.ptr.i34)
56-
%ptridx.i.i.i = getelementptr inbounds i32, i32 addrspace(1)* %add.ptr.i, i64 %4
57-
%ptridx.ascast.i.i.i = addrspacecast i32 addrspace(1)* %ptridx.i.i.i to i32 addrspace(4)*
58-
store i32 %call3.i.i.i.i, i32 addrspace(4)* %ptridx.ascast.i.i.i, align 4, !tbaa !14
41+
%5 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", ptr %_arg_9, i64 0, i32 0, i32 0, i64 0
42+
%6 = addrspacecast i64* %5 to ptr addrspace(4)
43+
%7 = load i64, ptr addrspace(4) %6, align 8
44+
%add.ptr.i = getelementptr inbounds i32, ptr addrspace(1) %_arg_5, i64 %7
45+
%8 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", ptr %_arg_4, i64 0, i32 0, i32 0, i64 0
46+
%9 = addrspacecast i64* %8 to ptr addrspace(4)
47+
%10 = load i64, ptr addrspace(4) %9, align 8
48+
%add.ptr.i34 = getelementptr inbounds i32, ptr addrspace(1) %_arg_1, i64 %10
49+
; CHECK: [[ARG_ASCAST:%[0-9a-zA-Z._]+]] = addrspacecast ptr addrspace(1) %[[ATOMIC_ARG_1:[0-9a-zA-Z._]+]] to ptr addrspace(4)
50+
; CHECK-NEXT: call void @__itt_offload_atomic_op_start(ptr addrspace(4) [[ARG_ASCAST]], i32 0, i32 0)
51+
; CHECK-NEXT: {{.*}}__spirv_AtomicLoad{{.*}}(ptr addrspace(1) %[[ATOMIC_ARG_1]],{{.*}}, i32 896
52+
; CHECK-NEXT: [[ARG_ASCAST:%[0-9a-zA-Z._]+]] = addrspacecast ptr addrspace(1) %[[ATOMIC_ARG_1]] to ptr addrspace(4)
53+
; CHECK-NEXT: call void @__itt_offload_atomic_op_finish(ptr addrspace(4) [[ARG_ASCAST]], i32 0, i32 0)
54+
%call3.i.i.i.i = tail call spir_func i32 @_Z18__spirv_AtomicLoadPU3AS1KiN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE(ptr addrspace(1) %add.ptr.i34, i32 1, i32 896) #2
55+
call spir_func void @__synthetic_spir_fun_call(ptr addrspace(1) %add.ptr.i34)
56+
%ptridx.i.i.i = getelementptr inbounds i32, ptr addrspace(1) %add.ptr.i, i64 %4
57+
%ptridx.ascast.i.i.i = addrspacecast ptr addrspace(1) %ptridx.i.i.i to ptr addrspace(4)
58+
store i32 %call3.i.i.i.i, ptr addrspace(4) %ptridx.ascast.i.i.i, align 4, !tbaa !14
5959
br label %_ZZN2cl4sycl7handler24parallel_for_lambda_implI11load_kernelIiEZZ9load_testIiEvNS0_5queueEmENKUlRS1_E_clES7_EUlNS0_4itemILi1ELb1EEEE_Li1EEEvNS0_5rangeIXT1_EEET0_ENKUlSA_E_clESA_.exit
6060

6161
_ZZN2cl4sycl7handler24parallel_for_lambda_implI11load_kernelIiEZZ9load_testIiEvNS0_5queueEmENKUlRS1_E_clES7_EUlNS0_4itemILi1ELb1EEEE_Li1EEEvNS0_5rangeIXT1_EEET0_ENKUlSA_E_clESA_.exit: ; preds = %entry, %if.end.i
@@ -64,56 +64,56 @@ _ZZN2cl4sycl7handler24parallel_for_lambda_implI11load_kernelIiEZZ9load_testIiEvN
6464
ret void
6565
}
6666

67-
define weak_odr dso_local spir_func void @__synthetic_spir_fun_call(i32 addrspace(1)* %ptr) {
67+
define weak_odr dso_local spir_func void @__synthetic_spir_fun_call(ptr addrspace(1) %ptr) {
6868
entry:
69-
; CHECK-LABEL: spir_func void @__synthetic_spir_fun_call(i32 addrspace(1)* %{{.*}}) {
69+
; CHECK-LABEL: spir_func void @__synthetic_spir_fun_call(ptr addrspace(1) %ptr) {
7070
; CHECK-NEXT: entry:
71-
; CHECK-NEXT: [[ARG_ASCAST:%[0-9a-zA-Z._]+]] = addrspacecast i32 addrspace(1)* %[[ATOMIC_ARG_S:[0-9a-zA-Z._]+]] to i8 addrspace(4)*
72-
; CHECK-NEXT: call void @__itt_offload_atomic_op_start(i8 addrspace(4)* [[ARG_ASCAST]], i32 0, i32 0)
73-
; CHECK-NEXT: {{.*}}__spirv_AtomicLoad{{.*}}(i32 addrspace(1)* %[[ATOMIC_ARG_S]],{{.*}}, i32 896
74-
; CHECK-NEXT: [[ARG_ASCAST:%[0-9a-zA-Z._]+]] = addrspacecast i32 addrspace(1)* %[[ATOMIC_ARG_S]] to i8 addrspace(4)*
75-
; CHECK-NEXT: call void @__itt_offload_atomic_op_finish(i8 addrspace(4)* [[ARG_ASCAST]], i32 0, i32 0)
76-
call spir_func i32 @_Z18__spirv_AtomicLoadPU3AS1KiN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE(i32 addrspace(1)* %ptr, i32 1, i32 896) #2
71+
; CHECK-NEXT: [[ARG_ASCAST:%[0-9a-zA-Z._]+]] = addrspacecast ptr addrspace(1) %[[ATOMIC_ARG_S:[0-9a-zA-Z._]+]] to ptr addrspace(4)
72+
; CHECK-NEXT: call void @__itt_offload_atomic_op_start(ptr addrspace(4) [[ARG_ASCAST]], i32 0, i32 0)
73+
; CHECK-NEXT: {{.*}}__spirv_AtomicLoad{{.*}}(ptr addrspace(1) %[[ATOMIC_ARG_S]],{{.*}}, i32 896
74+
; CHECK-NEXT: [[ARG_ASCAST:%[0-9a-zA-Z._]+]] = addrspacecast ptr addrspace(1) %[[ATOMIC_ARG_S]] to ptr addrspace(4)
75+
; CHECK-NEXT: call void @__itt_offload_atomic_op_finish(ptr addrspace(4) [[ARG_ASCAST]], i32 0, i32 0)
76+
call spir_func i32 @_Z18__spirv_AtomicLoadPU3AS1KiN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE(ptr addrspace(1) %ptr, i32 1, i32 896) #2
7777
; CHECK-NOT: call void @__itt_offload_wi_finish_wrapper()
7878
ret void
7979
}
8080

8181
; Function Attrs: convergent
82-
declare dso_local spir_func i32 @_Z18__spirv_AtomicLoadPU3AS1KiN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE(i32 addrspace(1)*, i32, i32) local_unnamed_addr #1
82+
declare dso_local spir_func i32 @_Z18__spirv_AtomicLoadPU3AS1KiN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE(ptr addrspace(1), i32, i32) local_unnamed_addr #1
8383

8484
; Function Attrs: convergent norecurse
85-
define weak_odr dso_local spir_kernel void @_ZTS11load_kernelIiE(i32 addrspace(1)* %_arg_, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_1, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_3, i32 addrspace(1)* %_arg_4, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_6, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_7, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_8) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !18 {
85+
define weak_odr dso_local spir_kernel void @_ZTS11load_kernelIiE(ptr addrspace(1) %_arg_, ptr byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_1, ptr byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, ptr byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_3, ptr addrspace(1) %_arg_4, ptr byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_6, ptr byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_7, ptr byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_8) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !18 {
8686
entry:
8787
; CHECK-LABEL: _ZTS11load_kernelIiE(
8888
; CHECK-NEXT: entry:
8989
; CHECK-NEXT: call void @__itt_offload_wi_start_wrapper()
90-
%0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0
91-
%1 = addrspacecast i64* %0 to i64 addrspace(4)*
92-
%2 = load i64, i64 addrspace(4)* %1, align 8
93-
%add.ptr.i32 = getelementptr inbounds i32, i32 addrspace(1)* %_arg_, i64 %2
94-
%3 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_8, i64 0, i32 0, i32 0, i64 0
95-
%4 = addrspacecast i64* %3 to i64 addrspace(4)*
96-
%5 = load i64, i64 addrspace(4)* %4, align 8
97-
%add.ptr.i = getelementptr inbounds i32, i32 addrspace(1)* %_arg_4, i64 %5
98-
%6 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId to <3 x i64> addrspace(4)*), align 32, !noalias !19
90+
%0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", ptr %_arg_3, i64 0, i32 0, i32 0, i64 0
91+
%1 = addrspacecast i64* %0 to ptr addrspace(4)
92+
%2 = load i64, ptr addrspace(4) %1, align 8
93+
%add.ptr.i32 = getelementptr inbounds i32, ptr addrspace(1) %_arg_, i64 %2
94+
%3 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", ptr %_arg_8, i64 0, i32 0, i32 0, i64 0
95+
%4 = addrspacecast i64* %3 to ptr addrspace(4)
96+
%5 = load i64, ptr addrspace(4) %4, align 8
97+
%add.ptr.i = getelementptr inbounds i32, ptr addrspace(1) %_arg_4, i64 %5
98+
%6 = load <3 x i64>, ptr addrspace(4) addrspacecast (ptr addrspace(1) @__spirv_BuiltInGlobalInvocationId to ptr addrspace(4)), align 32, !noalias !19
9999
%7 = extractelement <3 x i64> %6, i64 0
100-
; CHECK: [[ARG_ASCAST:%[0-9a-zA-Z._]+]] = addrspacecast i32 addrspace(1)* %[[ATOMIC_ARG_2:[0-9a-zA-Z._]+]] to i8 addrspace(4)*
101-
; CHECK-NEXT: call void @__itt_offload_atomic_op_start(i8 addrspace(4)* [[ARG_ASCAST]], i32 0, i32 0)
102-
; CHECK-NEXT: {{.*}}__spirv_AtomicLoad{{.*}}(i32 addrspace(1)* %[[ATOMIC_ARG_2]],{{.*}}, i32 896)
103-
; CHECK-NEXT: [[ARG_ASCAST:%[0-9a-zA-Z._]+]] = addrspacecast i32 addrspace(1)* %[[ATOMIC_ARG_2]] to i8 addrspace(4)*
104-
; CHECK-NEXT: call void @__itt_offload_atomic_op_finish(i8 addrspace(4)* [[ARG_ASCAST]], i32 0, i32 0)
105-
%call3.i.i.i = tail call spir_func i32 @_Z18__spirv_AtomicLoadPU3AS1KiN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE(i32 addrspace(1)* %add.ptr.i32, i32 1, i32 896) #2
106-
%ptridx.i.i = getelementptr inbounds i32, i32 addrspace(1)* %add.ptr.i, i64 %7
107-
%ptridx.ascast.i.i = addrspacecast i32 addrspace(1)* %ptridx.i.i to i32 addrspace(4)*
108-
store i32 %call3.i.i.i, i32 addrspace(4)* %ptridx.ascast.i.i, align 4, !tbaa !14
100+
; CHECK: [[ARG_ASCAST:%[0-9a-zA-Z._]+]] = addrspacecast ptr addrspace(1) %[[ATOMIC_ARG_2:[0-9a-zA-Z._]+]] to ptr addrspace(4)
101+
; CHECK-NEXT: call void @__itt_offload_atomic_op_start(ptr addrspace(4) [[ARG_ASCAST]], i32 0, i32 0)
102+
; CHECK-NEXT: {{.*}}__spirv_AtomicLoad{{.*}}(ptr addrspace(1) %[[ATOMIC_ARG_2]],{{.*}}, i32 896)
103+
; CHECK-NEXT: [[ARG_ASCAST:%[0-9a-zA-Z._]+]] = addrspacecast ptr addrspace(1) %[[ATOMIC_ARG_2]] to ptr addrspace(4)
104+
; CHECK-NEXT: call void @__itt_offload_atomic_op_finish(ptr addrspace(4) [[ARG_ASCAST]], i32 0, i32 0)
105+
%call3.i.i.i = tail call spir_func i32 @_Z18__spirv_AtomicLoadPU3AS1KiN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE(ptr addrspace(1) %add.ptr.i32, i32 1, i32 896) #2
106+
%ptridx.i.i = getelementptr inbounds i32, ptr addrspace(1) %add.ptr.i, i64 %7
107+
%ptridx.ascast.i.i = addrspacecast ptr addrspace(1) %ptridx.i.i to ptr addrspace(4)
108+
store i32 %call3.i.i.i, ptr addrspace(4) %ptridx.ascast.i.i, align 4, !tbaa !14
109109
; CHECK: call void @__itt_offload_wi_finish_wrapper()
110110
; CHECK-NEXT: ret void
111111
ret void
112112
}
113113

114114
; CHECK: declare void @__itt_offload_wi_start_wrapper()
115-
; CHECK: declare void @__itt_offload_atomic_op_start(i8 addrspace(4)*, i32, i32)
116-
; CHECK: declare void @__itt_offload_atomic_op_finish(i8 addrspace(4)*, i32, i32)
115+
; CHECK: declare void @__itt_offload_atomic_op_start(ptr addrspace(4), i32, i32)
116+
; CHECK: declare void @__itt_offload_atomic_op_finish(ptr addrspace(4), i32, i32)
117117
; CHECK: declare void @__itt_offload_wi_finish_wrapper()
118118

119119
attributes #0 = { convergent norecurse "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="llvm-test-suite/SYCL/AtomicRef/load.cpp" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" }

0 commit comments

Comments
 (0)