378 lines · plain
1// RUN: mlir-translate --mlir-to-llvmir %s | FileCheck %s2 3// CHECK-LABEL: @nvvm_tcgen05_ld_16x64b4llvm.func @nvvm_tcgen05_ld_16x64b(5 %tmemAddr : !llvm.ptr<6>,6 %stv1 : i32,7 %stv2 : vector<2xi32>,8 %stv4 : vector<4xi32>,9 %stv8 : vector<8xi32>,10 %stv16 : vector<16xi32>,11 %stv32 : vector<32xi32>,12 %stv64 : vector<64xi32>,13 %stv128 : vector<128xi32>) {14 15// CHECK: call void @llvm.nvvm.tcgen05.st.16x64b.x1(ptr addrspace(6) {{%[0-9]+}}, i32 {{%[0-9]+}}, i1 false)16 nvvm.tcgen05.st %tmemAddr, %stv1 { shape = #nvvm.tcgen05_ldst_shape<shape_16x64b>, num=1:i32 } : i3217 18// CHECK: call void @llvm.nvvm.tcgen05.st.16x64b.x2(ptr addrspace(6) {{%[0-9]+}}, <2 x i32> {{%[0-9]+}}, i1 false)19 nvvm.tcgen05.st %tmemAddr, %stv2 { shape = #nvvm.tcgen05_ldst_shape<shape_16x64b>, num=2:i32 } : vector<2xi32>20 21// CHECK: call void @llvm.nvvm.tcgen05.st.16x64b.x4(ptr addrspace(6) {{%[0-9]+}}, <4 x i32> {{%[0-9]+}}, i1 false)22 nvvm.tcgen05.st %tmemAddr, %stv4 { shape = #nvvm.tcgen05_ldst_shape<shape_16x64b>, num=4:i32 } : vector<4xi32>23 24// CHECK: call void @llvm.nvvm.tcgen05.st.16x64b.x8(ptr addrspace(6) {{%[0-9]+}}, <8 x i32> {{%[0-9]+}}, i1 false)25 nvvm.tcgen05.st %tmemAddr, %stv8 { shape = #nvvm.tcgen05_ldst_shape<shape_16x64b>, num=8:i32 } : vector<8xi32>26 27// CHECK: call void @llvm.nvvm.tcgen05.st.16x64b.x16(ptr addrspace(6) {{%[0-9]+}}, <16 x i32> {{%[0-9]+}}, i1 false)28 nvvm.tcgen05.st %tmemAddr, %stv16 { shape = #nvvm.tcgen05_ldst_shape<shape_16x64b>, num=16:i32 } : vector<16xi32>29 30// CHECK: call void @llvm.nvvm.tcgen05.st.16x64b.x32(ptr addrspace(6) {{%[0-9]+}}, <32 x i32> {{%[0-9]+}}, i1 false)31 nvvm.tcgen05.st %tmemAddr, %stv32 { shape = #nvvm.tcgen05_ldst_shape<shape_16x64b>, num=32:i32 } : vector<32xi32>32 33// CHECK: call void @llvm.nvvm.tcgen05.st.16x64b.x64(ptr addrspace(6) {{%[0-9]+}}, <64 x i32> {{%[0-9]+}}, i1 false)34 nvvm.tcgen05.st %tmemAddr, %stv64 { shape = #nvvm.tcgen05_ldst_shape<shape_16x64b>, num=64:i32 } : vector<64xi32>35 36// CHECK: call void @llvm.nvvm.tcgen05.st.16x64b.x128(ptr addrspace(6) {{%[0-9]+}}, <128 x i32> {{%[0-9]+}}, i1 false)37 nvvm.tcgen05.st %tmemAddr, %stv128 { shape = #nvvm.tcgen05_ldst_shape<shape_16x64b>, num=128:i32 } : vector<128xi32>38 39 llvm.return40}41 42// CHECK-LABEL: @nvvm_tcgen05_ld_16x64b_pack43llvm.func @nvvm_tcgen05_ld_16x64b_pack(44 %tmemAddr : !llvm.ptr<6>,45 %stv1 : i32,46 %stv2 : vector<2xi32>,47 %stv4 : vector<4xi32>,48 %stv8 : vector<8xi32>,49 %stv16 : vector<16xi32>,50 %stv32 : vector<32xi32>,51 %stv64 : vector<64xi32>,52 %stv128 : vector<128xi32>) {53 54// CHECK: call void @llvm.nvvm.tcgen05.st.16x64b.x1(ptr addrspace(6) {{%[0-9]+}}, i32 {{%[0-9]+}}, i1 true)55 nvvm.tcgen05.st %tmemAddr, %stv1 unpack { shape = #nvvm.tcgen05_ldst_shape<shape_16x64b>, num=1:i32 } : i3256 57// CHECK: call void @llvm.nvvm.tcgen05.st.16x64b.x2(ptr addrspace(6) {{%[0-9]+}}, <2 x i32> {{%[0-9]+}}, i1 true)58 nvvm.tcgen05.st %tmemAddr, %stv2 unpack { shape = #nvvm.tcgen05_ldst_shape<shape_16x64b>, num=2:i32 } : vector<2xi32>59 60// CHECK: call void @llvm.nvvm.tcgen05.st.16x64b.x4(ptr addrspace(6) {{%[0-9]+}}, <4 x i32> {{%[0-9]+}}, i1 true)61 nvvm.tcgen05.st %tmemAddr, %stv4 unpack { shape = #nvvm.tcgen05_ldst_shape<shape_16x64b>, num=4:i32 } : vector<4xi32>62 63// CHECK: call void @llvm.nvvm.tcgen05.st.16x64b.x8(ptr addrspace(6) {{%[0-9]+}}, <8 x i32> {{%[0-9]+}}, i1 true)64 nvvm.tcgen05.st %tmemAddr, %stv8 unpack { shape = #nvvm.tcgen05_ldst_shape<shape_16x64b>, num=8:i32 } : vector<8xi32>65 66// CHECK: call void @llvm.nvvm.tcgen05.st.16x64b.x16(ptr addrspace(6) {{%[0-9]+}}, <16 x i32> {{%[0-9]+}}, i1 true)67 nvvm.tcgen05.st %tmemAddr, %stv16 unpack { shape = #nvvm.tcgen05_ldst_shape<shape_16x64b>, num=16:i32 } : vector<16xi32>68 69// CHECK: call void @llvm.nvvm.tcgen05.st.16x64b.x32(ptr addrspace(6) {{%[0-9]+}}, <32 x i32> {{%[0-9]+}}, i1 true)70 nvvm.tcgen05.st %tmemAddr, %stv32 unpack { shape = #nvvm.tcgen05_ldst_shape<shape_16x64b>, num=32:i32 } : vector<32xi32>71 72// CHECK: call void @llvm.nvvm.tcgen05.st.16x64b.x64(ptr addrspace(6) {{%[0-9]+}}, <64 x i32> {{%[0-9]+}}, i1 true)73 nvvm.tcgen05.st %tmemAddr, %stv64 unpack { shape = #nvvm.tcgen05_ldst_shape<shape_16x64b>, num=64:i32 } : vector<64xi32>74 75// CHECK: call void @llvm.nvvm.tcgen05.st.16x64b.x128(ptr addrspace(6) {{%[0-9]+}}, <128 x i32> {{%[0-9]+}}, i1 true)76 nvvm.tcgen05.st %tmemAddr, %stv128 unpack { shape = #nvvm.tcgen05_ldst_shape<shape_16x64b>, num=128:i32 } : vector<128xi32>77 78 llvm.return79}80 81// CHECK-LABEL: @nvvm_tcgen05_ld_16x128b82llvm.func @nvvm_tcgen05_ld_16x128b(83 %tmemAddr : !llvm.ptr<6>,84 %stv1 : i32,85 %stv2 : vector<2xi32>,86 %stv4 : vector<4xi32>,87 %stv8 : vector<8xi32>,88 %stv16 : vector<16xi32>,89 %stv32 : vector<32xi32>,90 %stv64 : vector<64xi32>,91 %stv128 : vector<128xi32>) {92 93// CHECK: call void @llvm.nvvm.tcgen05.st.16x128b.x1(ptr addrspace(6) {{%[0-9]+}}, <2 x i32> {{%[0-9]+}}, i1 false)94 nvvm.tcgen05.st %tmemAddr, %stv2 { shape = #nvvm.tcgen05_ldst_shape<shape_16x128b>, num=1:i32 } : vector<2xi32>95 96// CHECK: call void @llvm.nvvm.tcgen05.st.16x128b.x2(ptr addrspace(6) {{%[0-9]+}}, <4 x i32> {{%[0-9]+}}, i1 false)97 nvvm.tcgen05.st %tmemAddr, %stv4 { shape = #nvvm.tcgen05_ldst_shape<shape_16x128b>, num=2:i32 } : vector<4xi32>98 99// CHECK: call void @llvm.nvvm.tcgen05.st.16x128b.x4(ptr addrspace(6) {{%[0-9]+}}, <8 x i32> {{%[0-9]+}}, i1 false)100 nvvm.tcgen05.st %tmemAddr, %stv8 { shape = #nvvm.tcgen05_ldst_shape<shape_16x128b>, num=4:i32 } : vector<8xi32>101 102// CHECK: call void @llvm.nvvm.tcgen05.st.16x128b.x8(ptr addrspace(6) {{%[0-9]+}}, <16 x i32> {{%[0-9]+}}, i1 false)103 nvvm.tcgen05.st %tmemAddr, %stv16 { shape = #nvvm.tcgen05_ldst_shape<shape_16x128b>, num=8:i32 } : vector<16xi32>104 105// CHECK: call void @llvm.nvvm.tcgen05.st.16x128b.x16(ptr addrspace(6) {{%[0-9]+}}, <32 x i32> {{%[0-9]+}}, i1 false)106 nvvm.tcgen05.st %tmemAddr, %stv32 { shape = #nvvm.tcgen05_ldst_shape<shape_16x128b>, num=16:i32 } : vector<32xi32>107 108// CHECK: call void @llvm.nvvm.tcgen05.st.16x128b.x32(ptr addrspace(6) {{%[0-9]+}}, <64 x i32> {{%[0-9]+}}, i1 false)109 nvvm.tcgen05.st %tmemAddr, %stv64 { shape = #nvvm.tcgen05_ldst_shape<shape_16x128b>, num=32:i32 } : vector<64xi32>110 111// CHECK: call void @llvm.nvvm.tcgen05.st.16x128b.x64(ptr addrspace(6) {{%[0-9]+}}, <128 x i32> {{%[0-9]+}}, i1 false)112 nvvm.tcgen05.st %tmemAddr, %stv128 { shape = #nvvm.tcgen05_ldst_shape<shape_16x128b>, num=64:i32 } : vector<128xi32>113 114 llvm.return115}116 117// CHECK-LABEL: @nvvm_tcgen05_ld_16x128b_pack118llvm.func @nvvm_tcgen05_ld_16x128b_pack(119 %tmemAddr : !llvm.ptr<6>,120 %stv1 : i32,121 %stv2 : vector<2xi32>,122 %stv4 : vector<4xi32>,123 %stv8 : vector<8xi32>,124 %stv16 : vector<16xi32>,125 %stv32 : vector<32xi32>,126 %stv64 : vector<64xi32>,127 %stv128 : vector<128xi32>) {128 129// CHECK: call void @llvm.nvvm.tcgen05.st.16x128b.x1(ptr addrspace(6) {{%[0-9]+}}, <2 x i32> {{%[0-9]+}}, i1 true)130 nvvm.tcgen05.st %tmemAddr, %stv2 unpack { shape = #nvvm.tcgen05_ldst_shape<shape_16x128b>, num=1:i32 } : vector<2xi32>131 132// CHECK: call void @llvm.nvvm.tcgen05.st.16x128b.x2(ptr addrspace(6) {{%[0-9]+}}, <4 x i32> {{%[0-9]+}}, i1 true)133 nvvm.tcgen05.st %tmemAddr, %stv4 unpack { shape = #nvvm.tcgen05_ldst_shape<shape_16x128b>, num=2:i32 } : vector<4xi32>134 135// CHECK: call void @llvm.nvvm.tcgen05.st.16x128b.x4(ptr addrspace(6) {{%[0-9]+}}, <8 x i32> {{%[0-9]+}}, i1 true)136 nvvm.tcgen05.st %tmemAddr, %stv8 unpack { shape = #nvvm.tcgen05_ldst_shape<shape_16x128b>, num=4:i32 } : vector<8xi32>137 138// CHECK: call void @llvm.nvvm.tcgen05.st.16x128b.x8(ptr addrspace(6) {{%[0-9]+}}, <16 x i32> {{%[0-9]+}}, i1 true)139 nvvm.tcgen05.st %tmemAddr, %stv16 unpack { shape = #nvvm.tcgen05_ldst_shape<shape_16x128b>, num=8:i32 } : vector<16xi32>140 141// CHECK: call void @llvm.nvvm.tcgen05.st.16x128b.x16(ptr addrspace(6) {{%[0-9]+}}, <32 x i32> {{%[0-9]+}}, i1 true)142 nvvm.tcgen05.st %tmemAddr, %stv32 unpack { shape = #nvvm.tcgen05_ldst_shape<shape_16x128b>, num=16:i32 } : vector<32xi32>143 144// CHECK: call void @llvm.nvvm.tcgen05.st.16x128b.x32(ptr addrspace(6) {{%[0-9]+}}, <64 x i32> {{%[0-9]+}}, i1 true)145 nvvm.tcgen05.st %tmemAddr, %stv64 unpack { shape = #nvvm.tcgen05_ldst_shape<shape_16x128b>, num=32:i32 } : vector<64xi32>146 147// CHECK: call void @llvm.nvvm.tcgen05.st.16x128b.x64(ptr addrspace(6) {{%[0-9]+}}, <128 x i32> {{%[0-9]+}}, i1 true)148 nvvm.tcgen05.st %tmemAddr, %stv128 unpack { shape = #nvvm.tcgen05_ldst_shape<shape_16x128b>, num=64:i32 } : vector<128xi32>149 150 llvm.return151}152 153// CHECK-LABEL: @nvvm_tcgen05_ld_16x256b154llvm.func @nvvm_tcgen05_ld_16x256b(155 %tmemAddr : !llvm.ptr<6>,156 %stv1 : i32,157 %stv2 : vector<2xi32>,158 %stv4 : vector<4xi32>,159 %stv8 : vector<8xi32>,160 %stv16 : vector<16xi32>,161 %stv32 : vector<32xi32>,162 %stv64 : vector<64xi32>,163 %stv128 : vector<128xi32>) {164 165// CHECK: call void @llvm.nvvm.tcgen05.st.16x256b.x1(ptr addrspace(6) {{%[0-9]+}}, <4 x i32> {{%[0-9]+}}, i1 false)166 nvvm.tcgen05.st %tmemAddr, %stv4 { shape = #nvvm.tcgen05_ldst_shape<shape_16x256b>, num=1:i32 } : vector<4xi32>167 168// CHECK: call void @llvm.nvvm.tcgen05.st.16x256b.x2(ptr addrspace(6) {{%[0-9]+}}, <8 x i32> {{%[0-9]+}}, i1 false)169 nvvm.tcgen05.st %tmemAddr, %stv8 { shape = #nvvm.tcgen05_ldst_shape<shape_16x256b>, num=2:i32 } : vector<8xi32>170 171// CHECK: call void @llvm.nvvm.tcgen05.st.16x256b.x4(ptr addrspace(6) {{%[0-9]+}}, <16 x i32> {{%[0-9]+}}, i1 false)172 nvvm.tcgen05.st %tmemAddr, %stv16 { shape = #nvvm.tcgen05_ldst_shape<shape_16x256b>, num=4:i32 } : vector<16xi32>173 174// CHECK: call void @llvm.nvvm.tcgen05.st.16x256b.x8(ptr addrspace(6) {{%[0-9]+}}, <32 x i32> {{%[0-9]+}}, i1 false)175 nvvm.tcgen05.st %tmemAddr, %stv32 { shape = #nvvm.tcgen05_ldst_shape<shape_16x256b>, num=8:i32 } : vector<32xi32>176 177// CHECK: call void @llvm.nvvm.tcgen05.st.16x256b.x16(ptr addrspace(6) {{%[0-9]+}}, <64 x i32> {{%[0-9]+}}, i1 false)178 nvvm.tcgen05.st %tmemAddr, %stv64 { shape = #nvvm.tcgen05_ldst_shape<shape_16x256b>, num=16:i32 } : vector<64xi32>179 180// CHECK: call void @llvm.nvvm.tcgen05.st.16x256b.x32(ptr addrspace(6) {{%[0-9]+}}, <128 x i32> {{%[0-9]+}}, i1 false)181 nvvm.tcgen05.st %tmemAddr, %stv128 { shape = #nvvm.tcgen05_ldst_shape<shape_16x256b>, num=32:i32 } : vector<128xi32>182 183 llvm.return184}185 186// CHECK-LABEL: @nvvm_tcgen05_ld_16x256b_pack187llvm.func @nvvm_tcgen05_ld_16x256b_pack(188 %tmemAddr : !llvm.ptr<6>,189 %stv1 : i32,190 %stv2 : vector<2xi32>,191 %stv4 : vector<4xi32>,192 %stv8 : vector<8xi32>,193 %stv16 : vector<16xi32>,194 %stv32 : vector<32xi32>,195 %stv64 : vector<64xi32>,196 %stv128 : vector<128xi32>) {197 198// CHECK: call void @llvm.nvvm.tcgen05.st.16x256b.x1(ptr addrspace(6) {{%[0-9]+}}, <4 x i32> {{%[0-9]+}}, i1 true)199 nvvm.tcgen05.st %tmemAddr, %stv4 unpack { shape = #nvvm.tcgen05_ldst_shape<shape_16x256b>, num=1:i32 } : vector<4xi32>200 201// CHECK: call void @llvm.nvvm.tcgen05.st.16x256b.x2(ptr addrspace(6) {{%[0-9]+}}, <8 x i32> {{%[0-9]+}}, i1 true)202 nvvm.tcgen05.st %tmemAddr, %stv8 unpack { shape = #nvvm.tcgen05_ldst_shape<shape_16x256b>, num=2:i32 } : vector<8xi32>203 204// CHECK: call void @llvm.nvvm.tcgen05.st.16x256b.x4(ptr addrspace(6) {{%[0-9]+}}, <16 x i32> {{%[0-9]+}}, i1 true)205 nvvm.tcgen05.st %tmemAddr, %stv16 unpack { shape = #nvvm.tcgen05_ldst_shape<shape_16x256b>, num=4:i32 } : vector<16xi32>206 207// CHECK: call void @llvm.nvvm.tcgen05.st.16x256b.x8(ptr addrspace(6) {{%[0-9]+}}, <32 x i32> {{%[0-9]+}}, i1 true)208 nvvm.tcgen05.st %tmemAddr, %stv32 unpack { shape = #nvvm.tcgen05_ldst_shape<shape_16x256b>, num=8:i32 } : vector<32xi32>209 210// CHECK: call void @llvm.nvvm.tcgen05.st.16x256b.x16(ptr addrspace(6) {{%[0-9]+}}, <64 x i32> {{%[0-9]+}}, i1 true)211 nvvm.tcgen05.st %tmemAddr, %stv64 unpack { shape = #nvvm.tcgen05_ldst_shape<shape_16x256b>, num=16:i32 } : vector<64xi32>212 213// CHECK: call void @llvm.nvvm.tcgen05.st.16x256b.x32(ptr addrspace(6) {{%[0-9]+}}, <128 x i32> {{%[0-9]+}}, i1 true)214 nvvm.tcgen05.st %tmemAddr, %stv128 unpack { shape = #nvvm.tcgen05_ldst_shape<shape_16x256b>, num=32:i32 } : vector<128xi32>215 216 llvm.return217}218 219// CHECK-LABEL: @nvvm_tcgen05_ld_32x32b220llvm.func @nvvm_tcgen05_ld_32x32b(221 %tmemAddr : !llvm.ptr<6>,222 %stv1 : i32,223 %stv2 : vector<2xi32>,224 %stv4 : vector<4xi32>,225 %stv8 : vector<8xi32>,226 %stv16 : vector<16xi32>,227 %stv32 : vector<32xi32>,228 %stv64 : vector<64xi32>,229 %stv128 : vector<128xi32>) {230 231// CHECK: call void @llvm.nvvm.tcgen05.st.32x32b.x1(ptr addrspace(6) {{%[0-9]+}}, i32 {{%[0-9]+}}, i1 false)232 nvvm.tcgen05.st %tmemAddr, %stv1 { shape = #nvvm.tcgen05_ldst_shape<shape_32x32b>, num=1:i32 } : i32233 234// CHECK: call void @llvm.nvvm.tcgen05.st.32x32b.x2(ptr addrspace(6) {{%[0-9]+}}, <2 x i32> {{%[0-9]+}}, i1 false)235 nvvm.tcgen05.st %tmemAddr, %stv2 { shape = #nvvm.tcgen05_ldst_shape<shape_32x32b>, num=2:i32 } : vector<2xi32>236 237// CHECK: call void @llvm.nvvm.tcgen05.st.32x32b.x4(ptr addrspace(6) {{%[0-9]+}}, <4 x i32> {{%[0-9]+}}, i1 false)238 nvvm.tcgen05.st %tmemAddr, %stv4 { shape = #nvvm.tcgen05_ldst_shape<shape_32x32b>, num=4:i32 } : vector<4xi32>239 240// CHECK: call void @llvm.nvvm.tcgen05.st.32x32b.x8(ptr addrspace(6) {{%[0-9]+}}, <8 x i32> {{%[0-9]+}}, i1 false)241 nvvm.tcgen05.st %tmemAddr, %stv8 { shape = #nvvm.tcgen05_ldst_shape<shape_32x32b>, num=8:i32 } : vector<8xi32>242 243// CHECK: call void @llvm.nvvm.tcgen05.st.32x32b.x16(ptr addrspace(6) {{%[0-9]+}}, <16 x i32> {{%[0-9]+}}, i1 false)244 nvvm.tcgen05.st %tmemAddr, %stv16 { shape = #nvvm.tcgen05_ldst_shape<shape_32x32b>, num=16:i32 } : vector<16xi32>245 246// CHECK: call void @llvm.nvvm.tcgen05.st.32x32b.x32(ptr addrspace(6) {{%[0-9]+}}, <32 x i32> {{%[0-9]+}}, i1 false)247 nvvm.tcgen05.st %tmemAddr, %stv32 { shape = #nvvm.tcgen05_ldst_shape<shape_32x32b>, num=32:i32 } : vector<32xi32>248 249// CHECK: call void @llvm.nvvm.tcgen05.st.32x32b.x64(ptr addrspace(6) {{%[0-9]+}}, <64 x i32> {{%[0-9]+}}, i1 false)250 nvvm.tcgen05.st %tmemAddr, %stv64 { shape = #nvvm.tcgen05_ldst_shape<shape_32x32b>, num=64:i32 } : vector<64xi32>251 252// CHECK: call void @llvm.nvvm.tcgen05.st.32x32b.x128(ptr addrspace(6) {{%[0-9]+}}, <128 x i32> {{%[0-9]+}}, i1 false)253 nvvm.tcgen05.st %tmemAddr, %stv128 { shape = #nvvm.tcgen05_ldst_shape<shape_32x32b>, num=128:i32 } : vector<128xi32>254 255 llvm.return256}257 258// CHECK-LABEL: @nvvm_tcgen05_ld_32x32b_pack259llvm.func @nvvm_tcgen05_ld_32x32b_pack(260 %tmemAddr : !llvm.ptr<6>,261 %stv1 : i32,262 %stv2 : vector<2xi32>,263 %stv4 : vector<4xi32>,264 %stv8 : vector<8xi32>,265 %stv16 : vector<16xi32>,266 %stv32 : vector<32xi32>,267 %stv64 : vector<64xi32>,268 %stv128 : vector<128xi32>) {269 270// CHECK: call void @llvm.nvvm.tcgen05.st.32x32b.x1(ptr addrspace(6) {{%[0-9]+}}, i32 {{%[0-9]+}}, i1 true)271 nvvm.tcgen05.st %tmemAddr, %stv1 unpack { shape = #nvvm.tcgen05_ldst_shape<shape_32x32b>, num=1:i32 } : i32272 273// CHECK: call void @llvm.nvvm.tcgen05.st.32x32b.x2(ptr addrspace(6) {{%[0-9]+}}, <2 x i32> {{%[0-9]+}}, i1 true)274 nvvm.tcgen05.st %tmemAddr, %stv2 unpack { shape = #nvvm.tcgen05_ldst_shape<shape_32x32b>, num=2:i32 } : vector<2xi32>275 276// CHECK: call void @llvm.nvvm.tcgen05.st.32x32b.x4(ptr addrspace(6) {{%[0-9]+}}, <4 x i32> {{%[0-9]+}}, i1 true)277 nvvm.tcgen05.st %tmemAddr, %stv4 unpack { shape = #nvvm.tcgen05_ldst_shape<shape_32x32b>, num=4:i32 } : vector<4xi32>278 279// CHECK: call void @llvm.nvvm.tcgen05.st.32x32b.x8(ptr addrspace(6) {{%[0-9]+}}, <8 x i32> {{%[0-9]+}}, i1 true)280 nvvm.tcgen05.st %tmemAddr, %stv8 unpack { shape = #nvvm.tcgen05_ldst_shape<shape_32x32b>, num=8:i32 } : vector<8xi32>281 282// CHECK: call void @llvm.nvvm.tcgen05.st.32x32b.x16(ptr addrspace(6) {{%[0-9]+}}, <16 x i32> {{%[0-9]+}}, i1 true)283 nvvm.tcgen05.st %tmemAddr, %stv16 unpack { shape = #nvvm.tcgen05_ldst_shape<shape_32x32b>, num=16:i32 } : vector<16xi32>284 285// CHECK: call void @llvm.nvvm.tcgen05.st.32x32b.x32(ptr addrspace(6) {{%[0-9]+}}, <32 x i32> {{%[0-9]+}}, i1 true)286 nvvm.tcgen05.st %tmemAddr, %stv32 unpack { shape = #nvvm.tcgen05_ldst_shape<shape_32x32b>, num=32:i32 } : vector<32xi32>287 288// CHECK: call void @llvm.nvvm.tcgen05.st.32x32b.x64(ptr addrspace(6) {{%[0-9]+}}, <64 x i32> {{%[0-9]+}}, i1 true)289 nvvm.tcgen05.st %tmemAddr, %stv64 unpack { shape = #nvvm.tcgen05_ldst_shape<shape_32x32b>, num=64:i32 } : vector<64xi32>290 291// CHECK: call void @llvm.nvvm.tcgen05.st.32x32b.x128(ptr addrspace(6) {{%[0-9]+}}, <128 x i32> {{%[0-9]+}}, i1 true)292 nvvm.tcgen05.st %tmemAddr, %stv128 unpack { shape = #nvvm.tcgen05_ldst_shape<shape_32x32b>, num=128:i32 } : vector<128xi32>293 294 llvm.return295}296 297// CHECK-LABEL: @nvvm_tcgen05_ld_16x32bx2298llvm.func @nvvm_tcgen05_ld_16x32bx2(299 %tmemAddr : !llvm.ptr<6>,300 %stv1 : i32,301 %stv2 : vector<2xi32>,302 %stv4 : vector<4xi32>,303 %stv8 : vector<8xi32>,304 %stv16 : vector<16xi32>,305 %stv32 : vector<32xi32>,306 %stv64 : vector<64xi32>,307 %stv128 : vector<128xi32>) {308 309 %offset = llvm.mlir.constant(2:i64) : i64310 311// CHECK: call void @llvm.nvvm.tcgen05.st.16x32bx2.x1(ptr addrspace(6) {{%[0-9]+}}, i64 2, i32 {{%[0-9]+}}, i1 false)312 nvvm.tcgen05.st %tmemAddr, %stv1, %offset { shape = #nvvm.tcgen05_ldst_shape<shape_16x32bx2>, num=1:i32 } : i32313 314// CHECK: call void @llvm.nvvm.tcgen05.st.16x32bx2.x2(ptr addrspace(6) {{%[0-9]+}}, i64 2, <2 x i32> {{%[0-9]+}}, i1 false)315 nvvm.tcgen05.st %tmemAddr, %stv2, %offset { shape = #nvvm.tcgen05_ldst_shape<shape_16x32bx2>, num=2:i32 } : vector<2xi32>316 317// CHECK: call void @llvm.nvvm.tcgen05.st.16x32bx2.x4(ptr addrspace(6) {{%[0-9]+}}, i64 2, <4 x i32> {{%[0-9]+}}, i1 false)318 nvvm.tcgen05.st %tmemAddr, %stv4, %offset { shape = #nvvm.tcgen05_ldst_shape<shape_16x32bx2>, num=4:i32 } : vector<4xi32>319 320// CHECK: call void @llvm.nvvm.tcgen05.st.16x32bx2.x8(ptr addrspace(6) {{%[0-9]+}}, i64 2, <8 x i32> {{%[0-9]+}}, i1 false)321 nvvm.tcgen05.st %tmemAddr, %stv8, %offset { shape = #nvvm.tcgen05_ldst_shape<shape_16x32bx2>, num=8:i32 } : vector<8xi32>322 323// CHECK: call void @llvm.nvvm.tcgen05.st.16x32bx2.x16(ptr addrspace(6) {{%[0-9]+}}, i64 2, <16 x i32> {{%[0-9]+}}, i1 false)324 nvvm.tcgen05.st %tmemAddr, %stv16, %offset { shape = #nvvm.tcgen05_ldst_shape<shape_16x32bx2>, num=16:i32 } : vector<16xi32>325 326// CHECK: call void @llvm.nvvm.tcgen05.st.16x32bx2.x32(ptr addrspace(6) {{%[0-9]+}}, i64 2, <32 x i32> {{%[0-9]+}}, i1 false)327 nvvm.tcgen05.st %tmemAddr, %stv32, %offset { shape = #nvvm.tcgen05_ldst_shape<shape_16x32bx2>, num=32:i32 } : vector<32xi32>328 329// CHECK: call void @llvm.nvvm.tcgen05.st.16x32bx2.x64(ptr addrspace(6) {{%[0-9]+}}, i64 2, <64 x i32> {{%[0-9]+}}, i1 false)330 nvvm.tcgen05.st %tmemAddr, %stv64, %offset { shape = #nvvm.tcgen05_ldst_shape<shape_16x32bx2>, num=64:i32 } : vector<64xi32>331 332// CHECK: call void @llvm.nvvm.tcgen05.st.16x32bx2.x128(ptr addrspace(6) {{%[0-9]+}}, i64 2, <128 x i32> {{%[0-9]+}}, i1 false)333 nvvm.tcgen05.st %tmemAddr, %stv128, %offset { shape = #nvvm.tcgen05_ldst_shape<shape_16x32bx2>, num=128:i32 } : vector<128xi32>334 335 llvm.return336}337 338// CHECK-LABEL: @nvvm_tcgen05_ld_16x32bx2_pack339llvm.func @nvvm_tcgen05_ld_16x32bx2_pack(340 %tmemAddr : !llvm.ptr<6>,341 %stv1 : i32,342 %stv2 : vector<2xi32>,343 %stv4 : vector<4xi32>,344 %stv8 : vector<8xi32>,345 %stv16 : vector<16xi32>,346 %stv32 : vector<32xi32>,347 %stv64 : vector<64xi32>,348 %stv128 : vector<128xi32>) {349 350 %offset = llvm.mlir.constant(2:i64) : i64351 352// CHECK: call void @llvm.nvvm.tcgen05.st.16x32bx2.x1(ptr addrspace(6) {{%[0-9]+}}, i64 2, i32 {{%[0-9]+}}, i1 true)353 nvvm.tcgen05.st %tmemAddr, %stv1, %offset unpack { shape = #nvvm.tcgen05_ldst_shape<shape_16x32bx2>, num=1:i32 } : i32354 355// CHECK: call void @llvm.nvvm.tcgen05.st.16x32bx2.x2(ptr addrspace(6) {{%[0-9]+}}, i64 2, <2 x i32> {{%[0-9]+}}, i1 true)356 nvvm.tcgen05.st %tmemAddr, %stv2, %offset unpack { shape = #nvvm.tcgen05_ldst_shape<shape_16x32bx2>, num=2:i32 } : vector<2xi32>357 358// CHECK: call void @llvm.nvvm.tcgen05.st.16x32bx2.x4(ptr addrspace(6) {{%[0-9]+}}, i64 2, <4 x i32> {{%[0-9]+}}, i1 true)359 nvvm.tcgen05.st %tmemAddr, %stv4, %offset unpack { shape = #nvvm.tcgen05_ldst_shape<shape_16x32bx2>, num=4:i32 } : vector<4xi32>360 361// CHECK: call void @llvm.nvvm.tcgen05.st.16x32bx2.x8(ptr addrspace(6) {{%[0-9]+}}, i64 2, <8 x i32> {{%[0-9]+}}, i1 true)362 nvvm.tcgen05.st %tmemAddr, %stv8, %offset unpack { shape = #nvvm.tcgen05_ldst_shape<shape_16x32bx2>, num=8:i32 } : vector<8xi32>363 364// CHECK: call void @llvm.nvvm.tcgen05.st.16x32bx2.x16(ptr addrspace(6) {{%[0-9]+}}, i64 2, <16 x i32> {{%[0-9]+}}, i1 true)365 nvvm.tcgen05.st %tmemAddr, %stv16, %offset unpack { shape = #nvvm.tcgen05_ldst_shape<shape_16x32bx2>, num=16:i32 } : vector<16xi32>366 367// CHECK: call void @llvm.nvvm.tcgen05.st.16x32bx2.x32(ptr addrspace(6) {{%[0-9]+}}, i64 2, <32 x i32> {{%[0-9]+}}, i1 true)368 nvvm.tcgen05.st %tmemAddr, %stv32, %offset unpack { shape = #nvvm.tcgen05_ldst_shape<shape_16x32bx2>, num=32:i32 } : vector<32xi32>369 370// CHECK: call void @llvm.nvvm.tcgen05.st.16x32bx2.x64(ptr addrspace(6) {{%[0-9]+}}, i64 2, <64 x i32> {{%[0-9]+}}, i1 true)371 nvvm.tcgen05.st %tmemAddr, %stv64, %offset unpack { shape = #nvvm.tcgen05_ldst_shape<shape_16x32bx2>, num=64:i32 } : vector<64xi32>372 373// CHECK: call void @llvm.nvvm.tcgen05.st.16x32bx2.x128(ptr addrspace(6) {{%[0-9]+}}, i64 2, <128 x i32> {{%[0-9]+}}, i1 true)374 nvvm.tcgen05.st %tmemAddr, %stv128, %offset unpack { shape = #nvvm.tcgen05_ldst_shape<shape_16x32bx2>, num=128:i32 } : vector<128xi32>375 376 llvm.return377}378