@@ -287,27 +287,27 @@ define void @nvvm_cp_async_bulk_intrinsics(ptr addrspace(3) %dst, ptr addrspace(
287287
288288; CHECK-LABEL: @nvvm_cp_async_bulk_tensor_g2s_im2col
289289define void @nvvm_cp_async_bulk_tensor_g2s_im2col (ptr addrspace (3 ) %d , ptr addrspace (3 ) %bar , ptr %tmap , i32 %d0 , i32 %d1 , i32 %d2 , i32 %d3 , i32 %d4 , i16 %im2col0 , i16 %im2col1 , i16 %im2col2 , i16 %mc , i64 %ch ) {
290- ; CHECK: call void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.3d(ptr addrspace(7) %1, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i16 %im2col0, i16 undef , i64 undef , i1 false, i1 false)
291- ; CHECK: call void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.4d(ptr addrspace(7) %2, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i16 %im2col0, i16 %im2col1, i16 undef , i64 undef , i1 false, i1 false)
292- ; CHECK: call void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.5d(ptr addrspace(7) %3, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, i16 %im2col0, i16 %im2col1, i16 %im2col2, i16 undef , i64 undef , i1 false, i1 false)
293- call void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.3d (ptr addrspace (3 ) %d , ptr addrspace (3 ) %bar , ptr %tmap , i32 %d0 , i32 %d1 , i32 %d2 , i16 %im2col0 , i16 undef , i64 undef , i1 0 , i1 0 )
294- call void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.4d (ptr addrspace (3 ) %d , ptr addrspace (3 ) %bar , ptr %tmap , i32 %d0 , i32 %d1 , i32 %d2 , i32 %d3 , i16 %im2col0 , i16 %im2col1 , i16 undef , i64 undef , i1 0 , i1 0 )
295- call void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.5d (ptr addrspace (3 ) %d , ptr addrspace (3 ) %bar , ptr %tmap , i32 %d0 , i32 %d1 , i32 %d2 , i32 %d3 , i32 %d4 , i16 %im2col0 , i16 %im2col1 , i16 %im2col2 , i16 undef , i64 undef , i1 0 , i1 0 )
290+ ; CHECK: call void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.3d(ptr addrspace(7) %1, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i16 %im2col0, i16 0 , i64 0 , i1 false, i1 false)
291+ ; CHECK: call void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.4d(ptr addrspace(7) %2, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i16 %im2col0, i16 %im2col1, i16 0 , i64 0 , i1 false, i1 false)
292+ ; CHECK: call void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.5d(ptr addrspace(7) %3, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, i16 %im2col0, i16 %im2col1, i16 %im2col2, i16 0 , i64 0 , i1 false, i1 false)
293+ call void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.3d (ptr addrspace (3 ) %d , ptr addrspace (3 ) %bar , ptr %tmap , i32 %d0 , i32 %d1 , i32 %d2 , i16 %im2col0 , i16 0 , i64 0 , i1 0 , i1 0 )
294+ call void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.4d (ptr addrspace (3 ) %d , ptr addrspace (3 ) %bar , ptr %tmap , i32 %d0 , i32 %d1 , i32 %d2 , i32 %d3 , i16 %im2col0 , i16 %im2col1 , i16 0 , i64 0 , i1 0 , i1 0 )
295+ call void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.5d (ptr addrspace (3 ) %d , ptr addrspace (3 ) %bar , ptr %tmap , i32 %d0 , i32 %d1 , i32 %d2 , i32 %d3 , i32 %d4 , i16 %im2col0 , i16 %im2col1 , i16 %im2col2 , i16 0 , i64 0 , i1 0 , i1 0 )
296296 ret void
297297}
298298
299299; CHECK-LABEL: @nvvm_cp_async_bulk_tensor_g2s_tile
300300define void @nvvm_cp_async_bulk_tensor_g2s_tile (ptr addrspace (3 ) %d , ptr addrspace (3 ) %bar , ptr %tmap , i32 %d0 , i32 %d1 , i32 %d2 , i32 %d3 , i32 %d4 , i16 %mc , i64 %ch ) {
301- ; CHECK: call void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.1d(ptr addrspace(7) %1, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i16 undef , i64 undef , i1 false, i1 false)
302- ; CHECK: call void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.2d(ptr addrspace(7) %2, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i32 %d1, i16 undef , i64 undef , i1 false, i1 false)
303- ; CHECK: call void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.3d(ptr addrspace(7) %3, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i16 undef , i64 undef , i1 false, i1 false)
304- ; CHECK: call void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.4d(ptr addrspace(7) %4, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i16 undef , i64 undef , i1 false, i1 false)
305- ; CHECK: call void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.5d(ptr addrspace(7) %5, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, i16 undef , i64 undef , i1 false, i1 false)
306- call void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.1d (ptr addrspace (3 ) %d , ptr addrspace (3 ) %bar , ptr %tmap , i32 %d0 , i16 undef , i64 undef , i1 0 , i1 0 )
307- call void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.2d (ptr addrspace (3 ) %d , ptr addrspace (3 ) %bar , ptr %tmap , i32 %d0 , i32 %d1 , i16 undef , i64 undef , i1 0 , i1 0 )
308- call void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.3d (ptr addrspace (3 ) %d , ptr addrspace (3 ) %bar , ptr %tmap , i32 %d0 , i32 %d1 , i32 %d2 , i16 undef , i64 undef , i1 0 , i1 0 )
309- call void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.4d (ptr addrspace (3 ) %d , ptr addrspace (3 ) %bar , ptr %tmap , i32 %d0 , i32 %d1 , i32 %d2 , i32 %d3 , i16 undef , i64 undef , i1 0 , i1 0 )
310- call void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.5d (ptr addrspace (3 ) %d , ptr addrspace (3 ) %bar , ptr %tmap , i32 %d0 , i32 %d1 , i32 %d2 , i32 %d3 , i32 %d4 , i16 undef , i64 undef , i1 0 , i1 0 )
301+ ; CHECK: call void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.1d(ptr addrspace(7) %1, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i16 0 , i64 0 , i1 false, i1 false)
302+ ; CHECK: call void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.2d(ptr addrspace(7) %2, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i32 %d1, i16 0 , i64 0 , i1 false, i1 false)
303+ ; CHECK: call void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.3d(ptr addrspace(7) %3, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i16 0 , i64 0 , i1 false, i1 false)
304+ ; CHECK: call void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.4d(ptr addrspace(7) %4, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i16 0 , i64 0 , i1 false, i1 false)
305+ ; CHECK: call void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.5d(ptr addrspace(7) %5, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, i16 0 , i64 0 , i1 false, i1 false)
306+ call void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.1d (ptr addrspace (3 ) %d , ptr addrspace (3 ) %bar , ptr %tmap , i32 %d0 , i16 0 , i64 0 , i1 0 , i1 0 )
307+ call void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.2d (ptr addrspace (3 ) %d , ptr addrspace (3 ) %bar , ptr %tmap , i32 %d0 , i32 %d1 , i16 0 , i64 0 , i1 0 , i1 0 )
308+ call void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.3d (ptr addrspace (3 ) %d , ptr addrspace (3 ) %bar , ptr %tmap , i32 %d0 , i32 %d1 , i32 %d2 , i16 0 , i64 0 , i1 0 , i1 0 )
309+ call void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.4d (ptr addrspace (3 ) %d , ptr addrspace (3 ) %bar , ptr %tmap , i32 %d0 , i32 %d1 , i32 %d2 , i32 %d3 , i16 0 , i64 0 , i1 0 , i1 0 )
310+ call void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.5d (ptr addrspace (3 ) %d , ptr addrspace (3 ) %bar , ptr %tmap , i32 %d0 , i32 %d1 , i32 %d2 , i32 %d3 , i32 %d4 , i16 0 , i64 0 , i1 0 , i1 0 )
311311 ret void
312312}
313313
0 commit comments