@@ -287,27 +287,27 @@ define void @nvvm_cp_async_bulk_intrinsics(ptr addrspace(3) %dst, ptr addrspace(
287
287
288
288
; CHECK-LABEL: @nvvm_cp_async_bulk_tensor_g2s_im2col
289
289
define 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 )
296
296
ret void
297
297
}
298
298
299
299
; CHECK-LABEL: @nvvm_cp_async_bulk_tensor_g2s_tile
300
300
define 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 )
311
311
ret void
312
312
}
313
313
0 commit comments