137 lines · plain
1// RUN: mlir-translate -mlir-to-llvmir %s | FileCheck %s2 3// CHECK-LABEL: @nvvm_tcgen05_cp_128x256b4llvm.func @nvvm_tcgen05_cp_128x256b(%taddr : !llvm.ptr<6>, %smem_desc : i64) {5 // CHECK: call void @llvm.nvvm.tcgen05.cp.128x256b.cg1(ptr addrspace(6) %{{.*}}, i64 %{{.*}})6 nvvm.tcgen05.cp %taddr, %smem_desc {shape = #nvvm.tcgen05_cp_shape<shape_128x256b>}7 8 // CHECK: call void @llvm.nvvm.tcgen05.cp.128x256b.cg2(ptr addrspace(6) %{{.*}}, i64 %{{.*}})9 nvvm.tcgen05.cp %taddr, %smem_desc {shape = #nvvm.tcgen05_cp_shape<shape_128x256b>, group = #nvvm.cta_group<cta_2>}10 11 // CHECK: call void @llvm.nvvm.tcgen05.cp.128x256b.b4x16_p64.cg2(ptr addrspace(6) %{{.*}}, i64 %{{.*}})12 nvvm.tcgen05.cp %taddr, %smem_desc {13 shape = #nvvm.tcgen05_cp_shape<shape_128x256b>,14 group = #nvvm.cta_group<cta_2>,15 srcFormat = #nvvm.tcgen05_cp_src_fmt<b4x16_p64>16 }17 // CHECK: call void @llvm.nvvm.tcgen05.cp.128x256b.b6x16_p32.cg2(ptr addrspace(6) %{{.*}}, i64 %{{.*}})18 nvvm.tcgen05.cp %taddr, %smem_desc {19 shape = #nvvm.tcgen05_cp_shape<shape_128x256b>,20 group = #nvvm.cta_group<cta_2>,21 srcFormat = #nvvm.tcgen05_cp_src_fmt<b6x16_p32>22 }23 llvm.return24}25 26// CHECK-LABEL: @nvvm_tcgen05_cp_4x256b27llvm.func @nvvm_tcgen05_cp_4x256b(%taddr : !llvm.ptr<6>, %smem_desc : i64) {28 // CHECK: call void @llvm.nvvm.tcgen05.cp.4x256b.cg1(ptr addrspace(6) %{{.*}}, i64 %{{.*}})29 nvvm.tcgen05.cp %taddr, %smem_desc {shape = #nvvm.tcgen05_cp_shape<shape_4x256b>}30 31 // CHECK: call void @llvm.nvvm.tcgen05.cp.4x256b.cg2(ptr addrspace(6) %{{.*}}, i64 %{{.*}})32 nvvm.tcgen05.cp %taddr, %smem_desc {shape = #nvvm.tcgen05_cp_shape<shape_4x256b>, group = #nvvm.cta_group<cta_2>}33 34 // CHECK: call void @llvm.nvvm.tcgen05.cp.4x256b.b4x16_p64.cg2(ptr addrspace(6) %{{.*}}, i64 %{{.*}})35 nvvm.tcgen05.cp %taddr, %smem_desc {36 shape = #nvvm.tcgen05_cp_shape<shape_4x256b>,37 group = #nvvm.cta_group<cta_2>,38 srcFormat = #nvvm.tcgen05_cp_src_fmt<b4x16_p64>39 }40 // CHECK: call void @llvm.nvvm.tcgen05.cp.4x256b.b6x16_p32.cg2(ptr addrspace(6) %{{.*}}, i64 %{{.*}})41 nvvm.tcgen05.cp %taddr, %smem_desc {42 shape = #nvvm.tcgen05_cp_shape<shape_4x256b>,43 group = #nvvm.cta_group<cta_2>,44 srcFormat = #nvvm.tcgen05_cp_src_fmt<b6x16_p32>45 }46 llvm.return47}48 49// CHECK-LABEL: @nvvm_tcgen05_cp_128x128b50llvm.func @nvvm_tcgen05_cp_128x128b(%taddr : !llvm.ptr<6>, %smem_desc : i64) {51 // CHECK: call void @llvm.nvvm.tcgen05.cp.128x128b.cg1(ptr addrspace(6) %{{.*}}, i64 %{{.*}})52 nvvm.tcgen05.cp %taddr, %smem_desc {shape = #nvvm.tcgen05_cp_shape<shape_128x128b>}53 54 // CHECK: call void @llvm.nvvm.tcgen05.cp.128x128b.cg2(ptr addrspace(6) %{{.*}}, i64 %{{.*}})55 nvvm.tcgen05.cp %taddr, %smem_desc {shape = #nvvm.tcgen05_cp_shape<shape_128x128b>, group = #nvvm.cta_group<cta_2>}56 57 // CHECK: call void @llvm.nvvm.tcgen05.cp.128x128b.b4x16_p64.cg2(ptr addrspace(6) %{{.*}}, i64 %{{.*}})58 nvvm.tcgen05.cp %taddr, %smem_desc {59 shape = #nvvm.tcgen05_cp_shape<shape_128x128b>,60 group = #nvvm.cta_group<cta_2>,61 srcFormat = #nvvm.tcgen05_cp_src_fmt<b4x16_p64>62 }63 // CHECK: call void @llvm.nvvm.tcgen05.cp.128x128b.b6x16_p32.cg2(ptr addrspace(6) %{{.*}}, i64 %{{.*}})64 nvvm.tcgen05.cp %taddr, %smem_desc {65 shape = #nvvm.tcgen05_cp_shape<shape_128x128b>,66 group = #nvvm.cta_group<cta_2>,67 srcFormat = #nvvm.tcgen05_cp_src_fmt<b6x16_p32>68 }69 llvm.return70}71 72// CHECK-LABEL: @nvvm_tcgen05_cp_64x128b73llvm.func @nvvm_tcgen05_cp_64x128b(%taddr : !llvm.ptr<6>, %smem_desc : i64) {74 // CHECK: call void @llvm.nvvm.tcgen05.cp.64x128b_warpx2_02_13.cg1(ptr addrspace(6) %{{.*}}, i64 %{{.*}})75 nvvm.tcgen05.cp %taddr, %smem_desc {76 shape = #nvvm.tcgen05_cp_shape<shape_64x128b>,77 multicast = #nvvm.tcgen05_cp_multicast<warpx2_02_13>78 }79 80 // CHECK: call void @llvm.nvvm.tcgen05.cp.64x128b_warpx2_02_13.cg2(ptr addrspace(6) %{{.*}}, i64 %{{.*}})81 nvvm.tcgen05.cp %taddr, %smem_desc {82 shape = #nvvm.tcgen05_cp_shape<shape_64x128b>,83 group = #nvvm.cta_group<cta_2>,84 multicast = #nvvm.tcgen05_cp_multicast<warpx2_02_13>85 }86 87 // CHECK: call void @llvm.nvvm.tcgen05.cp.64x128b_warpx2_02_13.b4x16_p64.cg1(ptr addrspace(6) %{{.*}}, i64 %{{.*}})88 nvvm.tcgen05.cp %taddr, %smem_desc {89 shape = #nvvm.tcgen05_cp_shape<shape_64x128b>,90 group = #nvvm.cta_group<cta_1>,91 multicast = #nvvm.tcgen05_cp_multicast<warpx2_02_13>,92 srcFormat = #nvvm.tcgen05_cp_src_fmt<b4x16_p64>93 }94 // CHECK: call void @llvm.nvvm.tcgen05.cp.64x128b_warpx2_01_23.b6x16_p32.cg2(ptr addrspace(6) %{{.*}}, i64 %{{.*}})95 nvvm.tcgen05.cp %taddr, %smem_desc {96 shape = #nvvm.tcgen05_cp_shape<shape_64x128b>,97 group = #nvvm.cta_group<cta_2>,98 multicast = #nvvm.tcgen05_cp_multicast<warpx2_01_23>,99 srcFormat = #nvvm.tcgen05_cp_src_fmt<b6x16_p32>100 }101 102 llvm.return103}104 105// CHECK-LABEL: @nvvm_tcgen05_cp_32x128b106llvm.func @nvvm_tcgen05_cp_32x128b(%taddr : !llvm.ptr<6>, %smem_desc : i64) {107 // CHECK: call void @llvm.nvvm.tcgen05.cp.32x128b_warpx4.cg1(ptr addrspace(6) %{{.*}}, i64 %{{.*}})108 nvvm.tcgen05.cp %taddr, %smem_desc {109 shape = #nvvm.tcgen05_cp_shape<shape_32x128b>,110 multicast = #nvvm.tcgen05_cp_multicast<warpx4>111 }112 113 // CHECK: call void @llvm.nvvm.tcgen05.cp.32x128b_warpx4.cg2(ptr addrspace(6) %{{.*}}, i64 %{{.*}})114 nvvm.tcgen05.cp %taddr, %smem_desc {115 shape = #nvvm.tcgen05_cp_shape<shape_32x128b>,116 group = #nvvm.cta_group<cta_2>,117 multicast = #nvvm.tcgen05_cp_multicast<warpx4>118 }119 120 // CHECK: call void @llvm.nvvm.tcgen05.cp.32x128b_warpx4.b4x16_p64.cg2(ptr addrspace(6) %{{.*}}, i64 %{{.*}})121 nvvm.tcgen05.cp %taddr, %smem_desc {122 shape = #nvvm.tcgen05_cp_shape<shape_32x128b>,123 group = #nvvm.cta_group<cta_2>,124 multicast = #nvvm.tcgen05_cp_multicast<warpx4>,125 srcFormat = #nvvm.tcgen05_cp_src_fmt<b4x16_p64>126 }127 // CHECK: call void @llvm.nvvm.tcgen05.cp.32x128b_warpx4.b6x16_p32.cg1(ptr addrspace(6) %{{.*}}, i64 %{{.*}})128 nvvm.tcgen05.cp %taddr, %smem_desc {129 shape = #nvvm.tcgen05_cp_shape<shape_32x128b>,130 group = #nvvm.cta_group<cta_1>,131 multicast = #nvvm.tcgen05_cp_multicast<warpx4>,132 srcFormat = #nvvm.tcgen05_cp_src_fmt<b6x16_p32>133 }134 135 llvm.return136}137