22
22
; CHECK-LLVM: %[[Apply:[%0-9a-z.]+]] = call spir_func target("spirv.CooperativeMatrixKHR", i16, 8, 16, 0, 3) @"_Z43__spirv_CooperativeMatrixApplyFunctionINTELPU3AS477class.sycl::_V1::ext::oneapi::experimental::matrix::helper::reference_wrapperPU3AS144__spirv_CooperativeMatrixKHR__short_8_16_0_3"(ptr addrspace(4) %ref.tmp.ascast.i21, target("spirv.CooperativeMatrixKHR", i16, 8, 16, 0, 3) %[[Mat]])
23
23
; CHECK-LLVM: call spir_func void @"_Z33__spirv_CooperativeMatrixStoreKHRPU3AS138class.sycl::_V1::ext::oneapi::bfloat16PU3AS144__spirv_CooperativeMatrixKHR__short_8_16_0_3liii"(ptr addrspace(1) %{{.*}}, target("spirv.CooperativeMatrixKHR", i16, 8, 16, 0, 3) %[[Apply]], i64 32, i32 0, i32 3, i32 0)
24
24
25
-
26
25
; ModuleID = 'matrix_apply.bc'
27
26
source_filename = "../llvm/sycl/test-e2e/Matrix/joint_matrix_apply_bf16.cpp"
28
27
target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
@@ -47,7 +46,6 @@ $_ZTSZZ17matrix_verify_addIN4sycl3_V13ext6oneapi8bfloat16ELm16ELm32EZ4mainEUlRS4
47
46
; Function Attrs: convergent norecurse nounwind
48
47
define weak_odr dso_local spir_kernel void @_ZTSZZ17matrix_verify_addIN4sycl3_V13ext6oneapi8bfloat16ELm16ELm32EZ4mainEUlRS4_E_EvNS1_5queueER10big_matrixIT_XT0_EXT1_EERNS1_8nd_rangeILi2EEEfOT2_ENKUlRNS1_7handlerEE_clESI_EUlNS1_7nd_itemILi2EEEE_ (ptr addrspace (1 ) noundef align 2 %_arg_accA , ptr noundef byval (%"class.sycl::_V1::range" ) align 8 %_arg_accA1 , ptr noundef byval (%"class.sycl::_V1::range" ) align 8 %_arg_accA2 , ptr noundef byval (%"class.sycl::_V1::id" ) align 8 %_arg_accA3 ) local_unnamed_addr {
49
48
entry:
50
- call spir_func void @__itt_offload_wi_start_wrapper ()
51
49
%ref.tmp.i20 = alloca %"class.sycl::_V1::ext::oneapi::experimental::matrix::helper::reference_wrapper" , align 8
52
50
%agg.tmp.i17 = alloca %"class.sycl::_V1::ext::oneapi::bfloat16" , align 2
53
51
%ref.tmp6.i = alloca float , align 4
@@ -85,14 +83,10 @@ entry:
85
83
%5 = load i64 , ptr addrspace (1 ) @__spirv_BuiltInLocalInvocationId , align 32
86
84
%ref.tmp6.ascast.i = addrspacecast ptr %ref.tmp6.i to ptr addrspace (4 )
87
85
%cmp.i11 = icmp ult i64 %2 , 2147483648
88
- tail call void @llvm.assume (i1 %cmp.i11 )
89
86
%cmp.i = icmp ult i64 %3 , 2147483648
90
- tail call void @llvm.assume (i1 %cmp.i )
91
87
%cmp.i15 = icmp ult i64 %4 , 2147483648
92
- tail call void @llvm.assume (i1 %cmp.i15 )
93
88
%sub.i = sub nsw i64 %2 , %4
94
89
%cmp.i12 = icmp ult i64 %5 , 2147483648
95
- tail call void @llvm.assume (i1 %cmp.i12 )
96
90
%sub5.i = sub nsw i64 %3 , %5
97
91
call void @llvm.lifetime.start.p0 (i64 4 , ptr nonnull %ref.tmp6.i )
98
92
store float 5 .000000e+00 , ptr %ref.tmp6.i , align 4
@@ -122,7 +116,6 @@ entry:
122
116
%add.ptr.i44 = getelementptr inbounds %"class.sycl::_V1::ext::oneapi::bfloat16" , ptr addrspace (1 ) %add.ptr.i43 , i64 %div14.i
123
117
call spir_func void @_Z33__spirv_CooperativeMatrixStoreKHRPU3AS4iPU3AS144__spirv_CooperativeMatrixKHR__uint_3_12_12_3ili (ptr addrspace (1 ) noundef %add.ptr.i44 , target ("spirv.CooperativeMatrixKHR" , i16 , 8 , 16 , 0 , 3 ) noundef %call.i22 , i64 noundef 32 , i32 noundef 0 , i32 noundef 3 , i32 noundef 0 )
124
118
call void @llvm.lifetime.end.p0 (i64 64 , ptr nonnull %__SYCLKernel )
125
- call spir_func void @__itt_offload_wi_finish_wrapper ()
126
119
ret void
127
120
}
128
121
@@ -132,9 +125,6 @@ declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture)
132
125
; Function Attrs: nocallback nofree nosync nounwind willreturn memory(argmem: readwrite)
133
126
declare void @llvm.lifetime.end.p0 (i64 immarg, ptr nocapture )
134
127
135
- ; Function Attrs: nocallback nofree nosync nounwind willreturn memory(inaccessiblemem: write)
136
- declare void @llvm.assume (i1 noundef)
137
-
138
128
; Function Attrs: convergent nounwind
139
129
declare dso_local spir_func noundef target ("spirv.CooperativeMatrixKHR" , i16 , 8 , 16 , 0 , 3 ) @_Z26__spirv_CompositeConstruct (ptr noundef byval (%"class.sycl::_V1::ext::oneapi::bfloat16" ) align 2 ) local_unnamed_addr
140
130
@@ -147,10 +137,6 @@ declare dso_local spir_func noundef target("spirv.CooperativeMatrixKHR", i16, 8,
147
137
; Function Attrs: convergent nounwind
148
138
declare dso_local spir_func void @_Z33__spirv_CooperativeMatrixStoreKHRPU3AS4iPU3AS144__spirv_CooperativeMatrixKHR__uint_3_12_12_3ili (ptr addrspace (1 ) noundef, target ("spirv.CooperativeMatrixKHR" , i16 , 8 , 16 , 0 , 3 ) noundef, i64 noundef, i32 noundef, i32 noundef, i32 noundef) local_unnamed_addr
149
139
150
- declare spir_func void @__itt_offload_wi_start_wrapper ()
151
-
152
- declare spir_func void @__itt_offload_wi_finish_wrapper ()
153
-
154
140
!llvm.module.flags = !{!0 , !1 }
155
141
!opencl.spir.version = !{!2 }
156
142
!spirv.Source = !{!3 }
0 commit comments