959 lines · plain
1// RUN: mlir-translate -mlir-to-llvmir %s -split-input-file --verify-diagnostics | FileCheck %s2 3// CHECK-LABEL: @nvvm_special_regs4llvm.func @nvvm_special_regs() -> i32 {5 // CHECK: %1 = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()6 %1 = nvvm.read.ptx.sreg.tid.x : i327 // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.tid.y()8 %2 = nvvm.read.ptx.sreg.tid.y : i329 // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.tid.z()10 %3 = nvvm.read.ptx.sreg.tid.z : i3211 // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()12 %4 = nvvm.read.ptx.sreg.ntid.x : i3213 // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ntid.y()14 %5 = nvvm.read.ptx.sreg.ntid.y : i3215 // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ntid.z()16 %6 = nvvm.read.ptx.sreg.ntid.z : i3217 // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()18 %7 = nvvm.read.ptx.sreg.ctaid.x : i3219 // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ctaid.y()20 %8 = nvvm.read.ptx.sreg.ctaid.y : i3221 // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.ctaid.z()22 %9 = nvvm.read.ptx.sreg.ctaid.z : i3223 // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.nctaid.x()24 %10 = nvvm.read.ptx.sreg.nctaid.x : i3225 // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.nctaid.y()26 %11 = nvvm.read.ptx.sreg.nctaid.y : i3227 // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.nctaid.z()28 %12 = nvvm.read.ptx.sreg.nctaid.z : i3229 // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.warpsize()30 %13 = nvvm.read.ptx.sreg.warpsize : i3231 // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.laneid()32 %14 = nvvm.read.ptx.sreg.laneid : i3233 // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.clusterid.x34 %15 = nvvm.read.ptx.sreg.clusterid.x : i3235 // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.clusterid.y36 %16 = nvvm.read.ptx.sreg.clusterid.y : i3237 // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.clusterid.z38 %17 = nvvm.read.ptx.sreg.clusterid.z : i3239 // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.nclusterid.x40 %18 = nvvm.read.ptx.sreg.nclusterid.x : i3241 // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.nclusterid.y42 %19 = nvvm.read.ptx.sreg.nclusterid.y : i3243 // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.nclusterid.z44 %20 = nvvm.read.ptx.sreg.nclusterid.z : i3245 // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.cluster.ctaid46 %21 = nvvm.read.ptx.sreg.cluster.ctaid.x : i3247 // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.cluster.ctaid48 %22 = nvvm.read.ptx.sreg.cluster.ctaid.y : i3249 // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.cluster.ctaid50 %23 = nvvm.read.ptx.sreg.cluster.ctaid.z : i3251 // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.cluster.nctaid52 %24 = nvvm.read.ptx.sreg.cluster.nctaid.x : i3253 // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.cluster.nctaid54 %25 = nvvm.read.ptx.sreg.cluster.nctaid.y : i3255 // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.cluster.nctaid56 %26 = nvvm.read.ptx.sreg.cluster.nctaid.z : i3257 // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.cluster.ctarank58 %27 = nvvm.read.ptx.sreg.cluster.ctarank : i3259 // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.cluster.nctarank60 %28 = nvvm.read.ptx.sreg.cluster.nctarank : i3261 // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.clock62 %29 = nvvm.read.ptx.sreg.clock : i3263 // CHECK: call i64 @llvm.nvvm.read.ptx.sreg.clock6464 %30 = nvvm.read.ptx.sreg.clock64 : i6465 // CHECK: call i64 @llvm.nvvm.read.ptx.sreg.globaltimer66 %31 = nvvm.read.ptx.sreg.globaltimer : i6467 // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.globaltimer.lo()68 %32 = nvvm.read.ptx.sreg.globaltimer.lo : i3269 // CHECK: %33 = call range(i32 0, 64) i32 @llvm.nvvm.read.ptx.sreg.tid.x()70 %33 = nvvm.read.ptx.sreg.tid.x range <i32, 0, 64> : i3271 // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.warpid72 %34 = nvvm.read.ptx.sreg.warpid : i3273 // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.nwarpid74 %35 = nvvm.read.ptx.sreg.nwarpid : i3275 // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.smid76 %36 = nvvm.read.ptx.sreg.smid : i3277 // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.nsmid78 %37 = nvvm.read.ptx.sreg.nsmid : i3279 // CHECK: call i32 @llvm.nvvm.read.ptx.sreg.gridid80 %38 = nvvm.read.ptx.sreg.gridid : i3281 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.envreg082 %39 = nvvm.read.ptx.sreg.envreg0 : i3283 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.envreg184 %40 = nvvm.read.ptx.sreg.envreg1 : i3285 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.envreg286 %41 = nvvm.read.ptx.sreg.envreg2 : i3287 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.envreg388 %42 = nvvm.read.ptx.sreg.envreg3 : i3289 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.envreg490 %43 = nvvm.read.ptx.sreg.envreg4 : i3291 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.envreg592 %44 = nvvm.read.ptx.sreg.envreg5 : i3293 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.envreg694 %45 = nvvm.read.ptx.sreg.envreg6 : i3295 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.envreg796 %46 = nvvm.read.ptx.sreg.envreg7 : i3297 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.envreg898 %47 = nvvm.read.ptx.sreg.envreg8 : i3299 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.envreg9100 %48 = nvvm.read.ptx.sreg.envreg9 : i32101 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.envreg10102 %49 = nvvm.read.ptx.sreg.envreg10 : i32103 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.envreg11104 %50 = nvvm.read.ptx.sreg.envreg11 : i32105 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.envreg12106 %51 = nvvm.read.ptx.sreg.envreg12 : i32107 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.envreg13108 %52 = nvvm.read.ptx.sreg.envreg13 : i32109 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.envreg14110 %53 = nvvm.read.ptx.sreg.envreg14 : i32111 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.envreg15112 %54 = nvvm.read.ptx.sreg.envreg15 : i32113 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.envreg16114 %55 = nvvm.read.ptx.sreg.envreg16 : i32115 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.envreg17116 %56 = nvvm.read.ptx.sreg.envreg17 : i32117 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.envreg18118 %57 = nvvm.read.ptx.sreg.envreg18 : i32119 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.envreg19120 %58 = nvvm.read.ptx.sreg.envreg19 : i32121 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.envreg20122 %59 = nvvm.read.ptx.sreg.envreg20 : i32123 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.envreg21124 %60 = nvvm.read.ptx.sreg.envreg21 : i32125 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.envreg22126 %61 = nvvm.read.ptx.sreg.envreg22 : i32127 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.envreg23128 %62 = nvvm.read.ptx.sreg.envreg23 : i32129 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.envreg24130 %63 = nvvm.read.ptx.sreg.envreg24 : i32131 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.envreg25132 %64 = nvvm.read.ptx.sreg.envreg25 : i32133 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.envreg26134 %65 = nvvm.read.ptx.sreg.envreg26 : i32135 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.envreg27136 %66 = nvvm.read.ptx.sreg.envreg27 : i32137 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.envreg28138 %67 = nvvm.read.ptx.sreg.envreg28 : i32139 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.envreg29140 %68 = nvvm.read.ptx.sreg.envreg29 : i32141 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.envreg30142 %69 = nvvm.read.ptx.sreg.envreg30 : i32143 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.envreg31144 %70 = nvvm.read.ptx.sreg.envreg31 : i32145 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.lanemask.eq146 %71 = nvvm.read.ptx.sreg.lanemask.eq : i32147 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.lanemask.le148 %72 = nvvm.read.ptx.sreg.lanemask.le : i32149 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.lanemask.lt150 %73 = nvvm.read.ptx.sreg.lanemask.lt : i32151 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.lanemask.ge152 %74 = nvvm.read.ptx.sreg.lanemask.ge : i32153 //CHECK: call i32 @llvm.nvvm.read.ptx.sreg.lanemask.gt154 %75 = nvvm.read.ptx.sreg.lanemask.gt : i32155 // CHECK: %76 = call range(i32 0, 0) i32 @llvm.nvvm.read.ptx.sreg.tid.x()156 %76 = nvvm.read.ptx.sreg.tid.x range <i32, 0, 0> : i32157 // CHECK: %77 = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()158 %77 = nvvm.read.ptx.sreg.tid.x range <i32, 4294967295, 4294967295> : i32159 llvm.return %1 : i32160}161 162// CHECK-LABEL: @nvvm_rcp163llvm.func @nvvm_rcp(%0: f32) -> f32 {164 // CHECK: call float @llvm.nvvm.rcp.approx.ftz.f165 %1 = nvvm.rcp.approx.ftz.f %0 : f32166 llvm.return %1 : f32167}168 169// CHECK-LABEL: @llvm_nvvm_cluster_arrive170llvm.func @llvm_nvvm_cluster_arrive() {171 // CHECK: call void @llvm.nvvm.barrier.cluster.arrive()172 nvvm.cluster.arrive173 // CHECK: call void @llvm.nvvm.barrier.cluster.arrive.aligned()174 nvvm.cluster.arrive {aligned}175 llvm.return176}177 178// CHECK-LABEL: @llvm_nvvm_cluster_arrive_relaxed179llvm.func @llvm_nvvm_cluster_arrive_relaxed() {180 // CHECK: call void @llvm.nvvm.barrier.cluster.arrive.relaxed()181 nvvm.cluster.arrive.relaxed182 // CHECK: call void @llvm.nvvm.barrier.cluster.arrive.relaxed.aligned()183 nvvm.cluster.arrive.relaxed {aligned}184 llvm.return185}186 187// CHECK-LABEL: @llvm_nvvm_cluster_wait188llvm.func @llvm_nvvm_cluster_wait() {189 // CHECK: call void @llvm.nvvm.barrier.cluster.wait()190 nvvm.cluster.wait191 // CHECK: call void @llvm.nvvm.barrier.cluster.wait.aligned()192 nvvm.cluster.wait {aligned}193 llvm.return194}195 196// CHECK-LABEL: @nvvm_shfl197llvm.func @nvvm_shfl(198 %0 : i32, %1 : i32, %2 : i32,199 %3 : i32, %4 : f32) -> i32 {200 // CHECK: call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}})201 %6 = nvvm.shfl.sync bfly %0, %3, %1, %2 : i32 -> i32202 // CHECK: call float @llvm.nvvm.shfl.sync.bfly.f32(i32 %{{.*}}, float %{{.*}}, i32 %{{.*}}, i32 %{{.*}})203 %7 = nvvm.shfl.sync bfly %0, %4, %1, %2 : f32 -> f32204 // CHECK: call i32 @llvm.nvvm.shfl.sync.up.i32(i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}})205 %8 = nvvm.shfl.sync up %0, %3, %1, %2 : i32 -> i32206 // CHECK: call float @llvm.nvvm.shfl.sync.up.f32(i32 %{{.*}}, float %{{.*}}, i32 %{{.*}}, i32 %{{.*}})207 %9 = nvvm.shfl.sync up %0, %4, %1, %2 : f32 -> f32208 // CHECK: call i32 @llvm.nvvm.shfl.sync.down.i32(i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}})209 %10 = nvvm.shfl.sync down %0, %3, %1, %2 : i32 -> i32210 // CHECK: call float @llvm.nvvm.shfl.sync.down.f32(i32 %{{.*}}, float %{{.*}}, i32 %{{.*}}, i32 %{{.*}})211 %11 = nvvm.shfl.sync down %0, %4, %1, %2 : f32 -> f32212 // CHECK: call i32 @llvm.nvvm.shfl.sync.idx.i32(i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}})213 %12 = nvvm.shfl.sync idx %0, %3, %1, %2 : i32 -> i32214 // CHECK: call float @llvm.nvvm.shfl.sync.idx.f32(i32 %{{.*}}, float %{{.*}}, i32 %{{.*}}, i32 %{{.*}})215 %13 = nvvm.shfl.sync idx %0, %4, %1, %2 : f32 -> f32216 llvm.return %6 : i32217}218 219// CHECK-LABEL: @nvvm_shfl_pred220llvm.func @nvvm_shfl_pred(221 %0 : i32, %1 : i32, %2 : i32,222 %3 : i32, %4 : f32) -> !llvm.struct<(i32, i1)> {223 // CHECK: call { i32, i1 } @llvm.nvvm.shfl.sync.bfly.i32p(i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}})224 %6 = nvvm.shfl.sync bfly %0, %3, %1, %2 {return_value_and_is_valid} : i32 -> !llvm.struct<(i32, i1)>225 // CHECK: call { float, i1 } @llvm.nvvm.shfl.sync.bfly.f32p(i32 %{{.*}}, float %{{.*}}, i32 %{{.*}}, i32 %{{.*}})226 %7 = nvvm.shfl.sync bfly %0, %4, %1, %2 {return_value_and_is_valid} : f32 -> !llvm.struct<(f32, i1)>227 // CHECK: call { i32, i1 } @llvm.nvvm.shfl.sync.up.i32p(i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}})228 %8 = nvvm.shfl.sync up %0, %3, %1, %2 {return_value_and_is_valid} : i32 -> !llvm.struct<(i32, i1)>229 // CHECK: call { float, i1 } @llvm.nvvm.shfl.sync.up.f32p(i32 %{{.*}}, float %{{.*}}, i32 %{{.*}}, i32 %{{.*}})230 %9 = nvvm.shfl.sync up %0, %4, %1, %2 {return_value_and_is_valid} : f32 -> !llvm.struct<(f32, i1)>231 // CHECK: call { i32, i1 } @llvm.nvvm.shfl.sync.down.i32p(i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}})232 %10 = nvvm.shfl.sync down %0, %3, %1, %2 {return_value_and_is_valid} : i32 -> !llvm.struct<(i32, i1)>233 // CHECK: call { float, i1 } @llvm.nvvm.shfl.sync.down.f32p(i32 %{{.*}}, float %{{.*}}, i32 %{{.*}}, i32 %{{.*}})234 %11 = nvvm.shfl.sync down %0, %4, %1, %2 {return_value_and_is_valid} : f32 -> !llvm.struct<(f32, i1)>235 // CHECK: call { i32, i1 } @llvm.nvvm.shfl.sync.idx.i32p(i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}})236 %12 = nvvm.shfl.sync idx %0, %3, %1, %2 {return_value_and_is_valid} : i32 -> !llvm.struct<(i32, i1)>237 // CHECK: call { float, i1 } @llvm.nvvm.shfl.sync.idx.f32p(i32 %{{.*}}, float %{{.*}}, i32 %{{.*}}, i32 %{{.*}})238 %13 = nvvm.shfl.sync idx %0, %4, %1, %2 {return_value_and_is_valid} : f32 -> !llvm.struct<(f32, i1)>239 llvm.return %6 : !llvm.struct<(i32, i1)>240}241 242// CHECK-LABEL: @nvvm_vote243llvm.func @nvvm_vote(%0 : i32, %1 : i1) -> i32 {244 // CHECK: call i32 @llvm.nvvm.vote.ballot.sync(i32 %{{.*}}, i1 %{{.*}})245 %3 = nvvm.vote.sync ballot %0, %1 -> i32246 // CHECK: call i1 @llvm.nvvm.vote.all.sync(i32 %{{.*}}, i1 %{{.*}})247 %4 = nvvm.vote.sync all %0, %1 -> i1248 // CHECK: call i1 @llvm.nvvm.vote.any.sync(i32 %{{.*}}, i1 %{{.*}})249 %5 = nvvm.vote.sync any %0, %1 -> i1250 // CHECK: call i1 @llvm.nvvm.vote.uni.sync(i32 %{{.*}}, i1 %{{.*}})251 %6 = nvvm.vote.sync uni %0, %1 -> i1252 llvm.return %3 : i32253}254 255// CHECK-LABEL: @nvvm_mma_mn8n8k4_row_col_f32_f32256llvm.func @nvvm_mma_mn8n8k4_row_col_f32_f32(%a0 : vector<2xf16>, %a1 : vector<2xf16>,257 %b0 : vector<2xf16>, %b1 : vector<2xf16>,258 %c0 : f32, %c1 : f32, %c2 : f32, %c3 : f32,259 %c4 : f32, %c5 : f32, %c6 : f32, %c7 : f32) -> !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)> {260 // CHECK: call { float, float, float, float, float, float, float, float } @llvm.nvvm.mma.m8n8k4.row.col.f32.f32261 %0 = nvvm.mma.sync A[%a0, %a1] B[%b0, %b1] C[%c0, %c1, %c2, %c3, %c4, %c5, %c6, %c7]262 {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>, shape = #nvvm.shape<m = 8, n = 8, k = 4>} : (vector<2xf16>, vector<2xf16>, f32) -> !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)>263 llvm.return %0 : !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)>264}265 266// CHECK-LABEL: @nvvm_mma_m16n8k16_f16_f16267llvm.func @nvvm_mma_m16n8k16_f16_f16(%a0 : vector<2xf16>, %a1 : vector<2xf16>,268 %a2 : vector<2xf16>, %a3 : vector<2xf16>,269 %b0 : vector<2xf16>, %b1 : vector<2xf16>,270 %c0 : vector<2xf16>, %c1 : vector<2xf16>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>)> {271 // CHECK: call { <2 x half>, <2 x half> } @llvm.nvvm.mma.m16n8k16.row.col.f16.f16272 %0 = nvvm.mma.sync A[ %a0, %a1, %a2, %a3 ] B[ %b0, %b1 ] C[ %c0, %c1 ]273 {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>, shape = #nvvm.shape<m = 16, n = 8, k = 16>}274 : (vector<2xf16>, vector<2xf16>, vector<2xf16>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>)>275 llvm.return %0 : !llvm.struct<(vector<2xf16>, vector<2xf16>)>276}277 278// CHECK-LABEL: @nvvm_mma_m16n8k16_bf16_bf16279llvm.func @nvvm_mma_m16n8k16_bf16_bf16(%a0 : i32, %a1 : i32, %a2 : i32, %a3 : i32,280 %b0 : i32, %b1 : i32,281 %c0 : f32, %c1 : f32, %c2 : f32, %c3 : f32) -> !llvm.struct<(f32, f32, f32, f32)> {282 // CHECK: call { float, float, float, float } @llvm.nvvm.mma.m16n8k16.row.col.bf16283 %0 = nvvm.mma.sync A[%a0, %a1, %a2, %a3] B[%b0, %b1] C[%c0, %c1, %c2, %c3]284 {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>,285 multiplicandAPtxType = #nvvm.mma_type<bf16>, multiplicandBPtxType = #nvvm.mma_type<bf16>,286 shape = #nvvm.shape<m = 16, n = 8, k = 16>} : (i32, i32, f32) -> !llvm.struct<(f32, f32, f32, f32)>287 llvm.return %0 : !llvm.struct<(f32, f32, f32, f32)>288}289 290// f32 return type, f32 accumulate type291// CHECK-LABEL: @nvvm_mma_m16n8k16_f32_f32292llvm.func @nvvm_mma_m16n8k16_f32_f32(%a0 : vector<2xf16>, %a1 : vector<2xf16>,293 %a2 : vector<2xf16>, %a3 : vector<2xf16>,294 %b0 : vector<2xf16>, %b1 : vector<2xf16>,295 %c0 : f32, %c1 : f32, %c2 : f32, %c3 : f32) -> !llvm.struct<(f32, f32, f32, f32)> {296 // CHECK: call { float, float, float, float } @llvm.nvvm.mma.m16n8k16.row.col.f32.f32297 %0 = nvvm.mma.sync A[%a0, %a1, %a2, %a3] B[%b0, %b1] C[%c0, %c1, %c2, %c3]298 {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>,299 shape = #nvvm.shape<m = 16, n = 8, k = 16>} : (vector<2xf16>, vector<2xf16>, f32) -> !llvm.struct<(f32, f32, f32, f32)>300 llvm.return %0 : !llvm.struct<(f32, f32, f32, f32)>301}302 303// CHECK-LABEL: @nvvm_mma_m16n8k16_s8_s8304llvm.func @nvvm_mma_m16n8k16_s8_s8(%a0 : i32, %a1 : i32,305 %b0 : i32,306 %c0 : i32, %c1 : i32, %c2 : i32, %c3 : i32) -> !llvm.struct<(i32, i32, i32, i32)> {307 // CHECK: call { i32, i32, i32, i32 } @llvm.nvvm.mma.m16n8k16.row.col.s8308 %0 = nvvm.mma.sync A[%a0, %a1] B[%b0] C[%c0, %c1, %c2, %c3]309 {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>,310 multiplicandAPtxType = #nvvm.mma_type<s8>, multiplicandBPtxType = #nvvm.mma_type<s8>,311 intOverflowBehavior=#nvvm.mma_int_overflow<wrapped>,312 shape = #nvvm.shape<m = 16, n = 8, k = 16>} : (i32, i32, i32) -> !llvm.struct<(i32,i32,i32,i32)>313 llvm.return %0 : !llvm.struct<(i32,i32,i32,i32)>314}315 316// CHECK-LABEL: @nvvm_mma_m16n8k16_s8_u8317llvm.func @nvvm_mma_m16n8k16_s8_u8(%a0 : i32, %a1 : i32,318 %b0 : i32,319 %c0 : i32, %c1 : i32, %c2 : i32, %c3 : i32) -> !llvm.struct<(i32, i32, i32, i32)> {320 // CHECK: call { i32, i32, i32, i32 } @llvm.nvvm.mma.m16n8k16.row.col.satfinite.s8.u8321 %0 = nvvm.mma.sync A[%a0, %a1] B[%b0] C[%c0, %c1, %c2, %c3]322 {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>,323 multiplicandAPtxType = #nvvm.mma_type<s8>, multiplicandBPtxType = #nvvm.mma_type<u8>,324 intOverflowBehavior=#nvvm.mma_int_overflow<satfinite>,325 shape = #nvvm.shape<m = 16, n = 8, k = 16>} : (i32, i32, i32) -> !llvm.struct<(i32,i32,i32,i32)>326 llvm.return %0 : !llvm.struct<(i32,i32,i32,i32)>327}328 329// CHECK-LABEL: @nvvm_mma_m16n8k128_b1_b1330llvm.func @nvvm_mma_m16n8k128_b1_b1(%a0 : i32, %a1 : i32,331 %b0 : i32,332 %c0 : i32, %c1 : i32, %c2 : i32, %c3 : i32) -> !llvm.struct<(i32,i32,i32,i32)> {333 // CHECK: call { i32, i32, i32, i32 } @llvm.nvvm.mma.xor.popc.m16n8k128.row.col.b1334 %0 = nvvm.mma.sync A[%a0, %a1] B[%b0] C[%c0, %c1, %c2, %c3]335 {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>,336 multiplicandAPtxType = #nvvm.mma_type<b1>, multiplicandBPtxType = #nvvm.mma_type<b1>,337 b1Op = #nvvm.mma_b1op<xor_popc>, shape = #nvvm.shape<m = 16, n = 8, k = 128>} : (i32, i32, i32) -> !llvm.struct<(i32,i32,i32,i32)>338 llvm.return %0 : !llvm.struct<(i32,i32,i32,i32)>339}340 341// CHECK-LABEL: @nvvm_mma_m16n8k32_s4_s4342llvm.func @nvvm_mma_m16n8k32_s4_s4(%a0 : i32, %a1 : i32,343 %b0 : i32,344 %c0 : i32, %c1 : i32, %c2 : i32, %c3 : i32) -> !llvm.struct<(i32,i32,i32,i32)> {345 // CHECK: call { i32, i32, i32, i32 } @llvm.nvvm.mma.m16n8k32.row.col.satfinite.s4346 %0 = nvvm.mma.sync A[%a0, %a1] B[%b0] C[%c0, %c1, %c2, %c3]347 {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>,348 multiplicandAPtxType = #nvvm.mma_type<s4>, multiplicandBPtxType = #nvvm.mma_type<s4>,349 intOverflowBehavior=#nvvm.mma_int_overflow<satfinite>,350 shape = #nvvm.shape<m = 16, n = 8, k = 32>} : (i32, i32, i32) -> !llvm.struct<(i32,i32,i32,i32)>351 llvm.return %0 : !llvm.struct<(i32,i32,i32,i32)>352}353 354// CHECK-LABEL: @nvvm_mma_m8n8k4_f64_f64355llvm.func @nvvm_mma_m8n8k4_f64_f64(%a0 : f64,356 %b0 : f64,357 %c0 : f64, %c1 : f64) -> !llvm.struct<(f64, f64)> {358 // CHECK: call { double, double } @llvm.nvvm.mma.m8n8k4.row.col.f64359 %0 = nvvm.mma.sync A[%a0] B[%b0] C[%c0, %c1]360 {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>,361 shape = #nvvm.shape<m = 8, n = 8, k = 4>} : (f64, f64, f64) -> !llvm.struct<(f64, f64)>362 llvm.return %0 : !llvm.struct<(f64, f64)>363}364 365// CHECK-LABEL: @nvvm_mma_m16n8k4_tf32_f32366llvm.func @nvvm_mma_m16n8k4_tf32_f32(%a0 : i32, %a1 : i32,367 %b0 : i32,368 %c0 : f32, %c1 : f32, %c2 : f32, %c3 : f32) -> !llvm.struct<(f32, f32, f32, f32)> {369 // CHECK: call { float, float, float, float } @llvm.nvvm.mma.m16n8k4.row.col.tf32370 %0 = nvvm.mma.sync A[%a0, %a1] B[%b0] C[%c0, %c1, %c2, %c3]371 {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>,372 multiplicandAPtxType = #nvvm.mma_type<tf32>, multiplicandBPtxType = #nvvm.mma_type<tf32>,373 shape = #nvvm.shape<m = 16, n = 8, k = 4>} : (i32, i32, f32) -> !llvm.struct<(f32, f32, f32, f32)>374 llvm.return %0 : !llvm.struct<(f32, f32, f32, f32)>375}376 377// The test below checks the correct mapping of the nvvm.wmma.*.load.* op to the correct intrinsic378// in the LLVM NVPTX backend.379// CHECK-LABEL: @gpu_wmma_load_op380llvm.func @gpu_wmma_load_op(%arg0: !llvm.ptr<3>, %arg1: i32) {381 // CHECK: call { <2 x half>, <2 x half>, <2 x half>, <2 x half>, <2 x half>, <2 x half>, <2 x half>, <2 x half> } @llvm.nvvm.wmma.m16n16k16.load.a.row.stride.f16.p3(ptr addrspace(3) %{{.*}}, i32 %{{.*}})382 %0 = nvvm.wmma.load %arg0, %arg1383 {eltype = #nvvm.mma_type<f16>, frag = #nvvm.mma_frag<a>, k = 16 : i32, layout = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}384 : (!llvm.ptr<3>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>385 386 llvm.return387}388 389// The test below checks the correct mapping of the nvvm.wmma.*.store.* op to the correct intrinsic390// in the LLVM NVPTX backend.391// CHECK-LABEL: @gpu_wmma_store_op392llvm.func @gpu_wmma_store_op(%arg0: !llvm.ptr<3>, %arg1: i32,393 %arg2: vector<2 x f16>, %arg3: vector<2 x f16>,394 %arg4: vector<2 xf16>, %arg5: vector<2 x f16>) {395 // CHECK: call void @llvm.nvvm.wmma.m16n16k16.store.d.row.stride.f16.p3(ptr addrspace(3) %{{.*}}, <2 x half> {{.*}}, <2 x half> %{{.*}}, <2 x half> %{{.*}}, <2 x half> %{{.*}}, i32 %{{.*}})396 nvvm.wmma.store %arg0, %arg1, %arg2, %arg3, %arg4, %arg5397 {eltype = #nvvm.mma_type<f16>, k = 16 : i32, layout = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}398 : !llvm.ptr<3>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>399 llvm.return400}401 402// The test below checks the correct mapping of the nvvm.wmma.*.mma.* op to the correct intrinsic403// in the LLVM NVPTX backend.404// CHECK-LABEL: @gpu_wmma_mma_op405llvm.func @gpu_wmma_mma_op(%arg0: vector<2 x f16>, %arg1: vector<2 x f16>,406 %arg2: vector<2 x f16>, %arg3: vector<2 x f16>,407 %arg4: vector<2 x f16>, %arg5: vector<2 x f16>,408 %arg6: vector<2 x f16>, %arg7: vector<2 x f16>,409 %arg8: vector<2 x f16>, %arg9: vector<2 x f16>,410 %arg10: vector<2 x f16>, %arg11: vector<2 x f16>,411 %arg12: vector<2 x f16>, %arg13: vector<2 x f16>,412 %arg14: vector<2 x f16>, %arg15: vector<2 x f16>,413 %arg16: vector<2 x f16>, %arg17: vector<2 x f16>,414 %arg18: vector<2 x f16>, %arg19: vector<2 x f16>) {415 // CHECK: call { <2 x half>, <2 x half>, <2 x half>, <2 x half> } @llvm.nvvm.wmma.m16n16k16.mma.row.row.f16.f16(<2 x half> {{.*}}, <2 x half> {{.*}}, <2 x half> {{.*}}, <2 x half> {{.*}}, <2 x half> {{.*}}, <2 x half> {{.*}}, <2 x half> {{.*}}, <2 x half> {{.*}}, <2 x half> {{.*}}, <2 x half> {{.*}}, <2 x half> {{.*}}, <2 x half> {{.*}}, <2 x half> {{.*}}, <2 x half> {{.*}}, <2 x half> {{.*}}, <2 x half> {{.*}}, <2 x half> {{.*}}, <2 x half> {{.*}}, <2 x half> {{.*}}, <2 x half> {{.*}})416 %0 = nvvm.wmma.mma %arg0, %arg1, %arg2, %arg3, %arg4, %arg5, %arg6, %arg7, %arg8, %arg9, %arg10, %arg11, %arg12, %arg13, %arg14, %arg15, %arg16, %arg17, %arg18, %arg19417 {eltypeA = #nvvm.mma_type<f16>, eltypeB = #nvvm.mma_type<f16>, k = 16 : i32, layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}418 : (vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>,419 vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>,420 vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>,421 vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>)422 -> !llvm.struct<(vector<2 x f16>, vector<2 x f16>, vector<2 x f16>, vector<2 x f16>)>423 llvm.return424}425 426// CHECK-LABEL: @nvvm_wmma_load_tf32427llvm.func @nvvm_wmma_load_tf32(%arg0: !llvm.ptr, %arg1 : i32) {428 // CHECK: call { i32, i32, i32, i32 } @llvm.nvvm.wmma.m16n16k8.load.a.row.stride.tf32.p0(ptr %{{.*}}, i32 %{{.*}})429 %0 = nvvm.wmma.load %arg0, %arg1430 {eltype = #nvvm.mma_type<tf32>, frag = #nvvm.mma_frag<a>, k = 8 : i32, layout = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}431 : (!llvm.ptr) -> !llvm.struct<(i32, i32, i32, i32)>432 llvm.return433}434 435// CHECK-LABEL: @nvvm_wmma_mma436llvm.func @nvvm_wmma_mma(%0 : i32, %1 : i32, %2 : i32, %3 : i32, %4 : i32, %5 : i32,437 %6 : i32, %7 : i32, %8 : f32, %9 : f32, %10 : f32,438 %11 : f32, %12 : f32, %13 : f32, %14 : f32, %15 : f32) {439 // CHECK: { float, float, float, float, float, float, float, float } @llvm.nvvm.wmma.m16n16k8.mma.row.row.tf32(i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, float %{{.*}}, float %{{.*}}, float %{{.*}}, float %{{.*}}, float %{{.*}}, float %{{.*}}, float %{{.*}}, float %{{.*}})440 %r = nvvm.wmma.mma %0, %1, %2, %3, %4, %5, %6, %7, %8, %9, %10, %11, %12, %13, %14, %15441 {eltypeA = #nvvm.mma_type<tf32>, eltypeB = #nvvm.mma_type<f32>, k = 8 : i32, layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}442 : (i32, i32, i32, i32, i32, i32, i32, i32, f32, f32, f32, f32, f32, f32, f32, f32)443 -> !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)>444 llvm.return445}446 447// CHECK-LABEL: @nvvm_wmma_load_a_f64448llvm.func @nvvm_wmma_load_a_f64(%arg0: !llvm.ptr, %arg1 : i32) {449 // CHECK: call double @llvm.nvvm.wmma.m8n8k4.load.a.row.stride.f64.p0(ptr %{{.*}}, i32 %{{.*}})450 %0 = nvvm.wmma.load %arg0, %arg1451 {eltype = #nvvm.mma_type<f64>, frag = #nvvm.mma_frag<a>, k = 4 : i32, layout = #nvvm.mma_layout<row>, m = 8 : i32, n = 8 : i32}452 : (!llvm.ptr) -> f64453 llvm.return454}455 456// CHECK-LABEL: @nvvm_wmma_load_c_f64457llvm.func @nvvm_wmma_load_c_f64(%arg0: !llvm.ptr, %arg1 : i32) {458 // CHECK: call { double, double } @llvm.nvvm.wmma.m8n8k4.load.c.row.stride.f64.p0(ptr %{{.*}}, i32 %{{.*}})459 %0 = nvvm.wmma.load %arg0, %arg1460 {eltype = #nvvm.mma_type<f64>, frag = #nvvm.mma_frag<c>, k = 4 : i32, layout = #nvvm.mma_layout<row>, m = 8 : i32, n = 8 : i32}461 : (!llvm.ptr) -> !llvm.struct<(f64, f64)>462 llvm.return463}464 465// CHECK-LABEL: @nvvm_wmma_mma_f64466llvm.func @nvvm_wmma_mma_f64(%0 : f64, %1 : f64, %2 : f64, %3 : f64) {467 // CHECK: { double, double } @llvm.nvvm.wmma.m8n8k4.mma.row.col.f64(double %{{.*}}, double %{{.*}}, double %{{.*}}, double %{{.*}})468 %r = nvvm.wmma.mma %0, %1, %2, %3469 {eltypeA = #nvvm.mma_type<f64>, eltypeB = #nvvm.mma_type<f64>, k = 4 : i32, layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>, m = 8 : i32, n = 8 : i32}470 : (f64, f64, f64, f64)471 -> !llvm.struct<(f64, f64)>472 llvm.return473}474 475// CHECK-LABEL: @nvvm_wmma_store_d_f64476llvm.func @nvvm_wmma_store_d_f64(%arg0: !llvm.ptr, %arg1 : i32, %arg2 : f64, %arg3 : f64) {477 // CHECK: call void @llvm.nvvm.wmma.m8n8k4.store.d.row.stride.f64.p0(ptr %{{.*}}, double %{{.*}}, double %{{.*}}, i32 %{{.*}})478 nvvm.wmma.store %arg0, %arg1, %arg2, %arg3479 {eltype = #nvvm.mma_type<f64>, k = 4 : i32, layout = #nvvm.mma_layout<row>, m = 8 : i32, n = 8 : i32}480 : !llvm.ptr, f64, f64481 llvm.return482}483 484// CHECK-LABEL: @cp_async485llvm.func @cp_async(%arg0: !llvm.ptr<3>, %arg1: !llvm.ptr<1>) {486 // CHECK: call void @llvm.nvvm.cp.async.ca.shared.global.4(ptr addrspace(3) %{{.*}}, ptr addrspace(1) %{{.*}})487 nvvm.cp.async.shared.global %arg0, %arg1, 4, cache = ca : !llvm.ptr<3>, !llvm.ptr<1>488 // CHECK: call void @llvm.nvvm.cp.async.ca.shared.global.8(ptr addrspace(3) %{{.*}}, ptr addrspace(1) %{{.*}})489 nvvm.cp.async.shared.global %arg0, %arg1, 8, cache = ca : !llvm.ptr<3>, !llvm.ptr<1>490 // CHECK: call void @llvm.nvvm.cp.async.ca.shared.global.16(ptr addrspace(3) %{{.*}}, ptr addrspace(1) %{{.*}})491 nvvm.cp.async.shared.global %arg0, %arg1, 16, cache = ca : !llvm.ptr<3>, !llvm.ptr<1>492 // CHECK: call void @llvm.nvvm.cp.async.cg.shared.global.16(ptr addrspace(3) %{{.*}}, ptr addrspace(1) %{{.*}})493 nvvm.cp.async.shared.global %arg0, %arg1, 16, cache = cg : !llvm.ptr<3>, !llvm.ptr<1>494 495 // CHECK: call void @llvm.nvvm.cp.async.commit.group()496 nvvm.cp.async.commit.group497 // CHECK: call void @llvm.nvvm.cp.async.wait.group(i32 0)498 nvvm.cp.async.wait.group 0499 llvm.return500}501 502// CHECK-LABEL: @async_cp_zfill503llvm.func @async_cp_zfill(%dst: !llvm.ptr<3>, %src: !llvm.ptr<1>, %cpSize: i32) {504 // CHECK: call void @llvm.nvvm.cp.async.ca.shared.global.4.s(ptr addrspace(3) %{{.*}}, ptr addrspace(1) %{{.*}}, i32 %{{.*}})505 nvvm.cp.async.shared.global %dst, %src, 4, cache = ca, %cpSize : !llvm.ptr<3>, !llvm.ptr<1>, i32506 // CHECK: call void @llvm.nvvm.cp.async.ca.shared.global.8.s(ptr addrspace(3) %{{.*}}, ptr addrspace(1) %{{.*}}, i32 %{{.*}})507 nvvm.cp.async.shared.global %dst, %src, 8, cache = ca, %cpSize : !llvm.ptr<3>, !llvm.ptr<1>, i32508 // CHECK: call void @llvm.nvvm.cp.async.ca.shared.global.16.s(ptr addrspace(3) %{{.*}}, ptr addrspace(1) %{{.*}}, i32 %{{.*}})509 nvvm.cp.async.shared.global %dst, %src, 16, cache = ca, %cpSize : !llvm.ptr<3>, !llvm.ptr<1>, i32510 // CHECK: call void @llvm.nvvm.cp.async.cg.shared.global.16.s(ptr addrspace(3) %{{.*}}, ptr addrspace(1) %{{.*}}, i32 %{{.*}})511 nvvm.cp.async.shared.global %dst, %src, 16, cache = cg, %cpSize : !llvm.ptr<3>, !llvm.ptr<1>, i32512 llvm.return513}514 515// CHECK-LABEL: @llvm_nvvm_setmaxregister516llvm.func @llvm_nvvm_setmaxregister() {517 // CHECK: call void @llvm.nvvm.setmaxnreg.inc.sync.aligned.u32(i32 256)518 nvvm.setmaxregister increase 256519 // CHECK: call void @llvm.nvvm.setmaxnreg.dec.sync.aligned.u32(i32 24)520 nvvm.setmaxregister decrease 24521 llvm.return522}523 524// CHECK-LABEL: @llvm_nvvm_cp_async_bulk_commit_group525llvm.func @llvm_nvvm_cp_async_bulk_commit_group() {526 // CHECK: call void @llvm.nvvm.cp.async.bulk.commit.group()527 nvvm.cp.async.bulk.commit.group528 llvm.return529}530 531// CHECK-LABEL: @llvm_nvvm_cp_async_bulk_wait_group532llvm.func @llvm_nvvm_cp_async_bulk_wait_group() {533 // CHECK: call void @llvm.nvvm.cp.async.bulk.wait.group(i32 0)534 nvvm.cp.async.bulk.wait_group 0535 // CHECK: call void @llvm.nvvm.cp.async.bulk.wait.group(i32 3)536 nvvm.cp.async.bulk.wait_group 3537 // CHECK: call void @llvm.nvvm.cp.async.bulk.wait.group.read(i32 0)538 nvvm.cp.async.bulk.wait_group 0 {read}539 // CHECK: call void @llvm.nvvm.cp.async.bulk.wait.group.read(i32 3)540 nvvm.cp.async.bulk.wait_group 3 {read}541 llvm.return542}543 544// CHECK-LABEL: @ld_matrix545llvm.func @ld_matrix(%arg0: !llvm.ptr<3>) {546 // CHECK: call i32 @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x1.b16.p3(ptr addrspace(3) %{{.*}})547 %l1 = nvvm.ldmatrix %arg0 {num = 1: i32, layout = #nvvm.mma_layout<row>, shape = #nvvm.ld_st_matrix_shape<m = 8, n = 8>, eltType = #nvvm.ld_st_matrix_elt_type<b16>} : (!llvm.ptr<3>) -> i32548 // CHECK: call { i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x2.b16.p3(ptr addrspace(3) %{{.*}})549 %l2 = nvvm.ldmatrix %arg0 {num = 2 : i32, layout = #nvvm.mma_layout<row>, shape =#nvvm.ld_st_matrix_shape<m = 8, n = 8>, eltType = #nvvm.ld_st_matrix_elt_type<b16>} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32)>550 // CHECK: call { i32, i32, i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x4.b16.p3(ptr addrspace(3) %{{.*}})551 %l4 = nvvm.ldmatrix %arg0 {num = 4 : i32, layout = #nvvm.mma_layout<row>, shape =#nvvm.ld_st_matrix_shape<m = 8, n = 8>, eltType = #nvvm.ld_st_matrix_elt_type<b16>} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32, i32, i32)>552 553 // CHECK: call i32 @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x1.trans.b16.p3(ptr addrspace(3) %{{.*}})554 %l1t = nvvm.ldmatrix %arg0 {num = 1: i32, layout = #nvvm.mma_layout<col>, shape =#nvvm.ld_st_matrix_shape<m = 8, n = 8>, eltType = #nvvm.ld_st_matrix_elt_type<b16>} : (!llvm.ptr<3>) -> i32555 // CHECK: call { i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x2.trans.b16.p3(ptr addrspace(3) %{{.*}})556 %l2t = nvvm.ldmatrix %arg0 {num = 2 : i32, layout = #nvvm.mma_layout<col>, shape =#nvvm.ld_st_matrix_shape<m = 8, n = 8>, eltType = #nvvm.ld_st_matrix_elt_type<b16>} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32)>557 // CHECK: call { i32, i32, i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x4.trans.b16.p3(ptr addrspace(3) %{{.*}})558 %l4t = nvvm.ldmatrix %arg0 {num = 4 : i32, layout = #nvvm.mma_layout<col>, shape = #nvvm.ld_st_matrix_shape<m = 8, n = 8>, eltType = #nvvm.ld_st_matrix_elt_type<b16>} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32, i32, i32)>559 560 // CHECK: call i32 @llvm.nvvm.ldmatrix.sync.aligned.m8n16.x1.b8x16.b6x16_p32.p3(ptr addrspace(3) %{{.*}})561 %m8n16_b6_l1 = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout<row>, shape =#nvvm.ld_st_matrix_shape<m = 8, n = 16>, eltType = #nvvm.ld_st_matrix_elt_type<b8x16.b6x16_p32>} : (!llvm.ptr<3>) -> i32562 // CHECK: call { i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n16.x2.b8x16.b6x16_p32.p3(ptr addrspace(3) %{{.*}})563 %m8n16_b6_l2 = nvvm.ldmatrix %arg0 {num = 2: i32, layout = #nvvm.mma_layout<row>, shape =#nvvm.ld_st_matrix_shape<m = 8, n = 16>, eltType = #nvvm.ld_st_matrix_elt_type<b8x16.b6x16_p32>} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32)>564 // CHECK: call { i32, i32, i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n16.x4.b8x16.b6x16_p32.p3(ptr addrspace(3) %{{.*}})565 %m8n16_b6_l4 = nvvm.ldmatrix %arg0{num = 4 : i32, layout = #nvvm.mma_layout<row>, shape = #nvvm.ld_st_matrix_shape<m = 8, n = 16>,eltType =#nvvm.ld_st_matrix_elt_type<b8x16.b6x16_p32>} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32, i32, i32)>566 567 // CHECK: call i32 @llvm.nvvm.ldmatrix.sync.aligned.m8n16.x1.b8x16.b4x16_p64.p3(ptr addrspace(3) %{{.*}})568 %m8n16_b4_l1 = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout<row>, shape =#nvvm.ld_st_matrix_shape<m = 8, n = 16>, eltType = #nvvm.ld_st_matrix_elt_type<b8x16.b4x16_p64>} : (!llvm.ptr<3>) -> i32569 // CHECK: call { i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n16.x2.b8x16.b4x16_p64.p3(ptr addrspace(3) %{{.*}})570 %m8n16_b4_l2 = nvvm.ldmatrix %arg0 {num = 2 : i32, layout = #nvvm.mma_layout<row>, shape = #nvvm.ld_st_matrix_shape<m = 8, n = 16>, eltType = #nvvm.ld_st_matrix_elt_type<b8x16.b4x16_p64>} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32)>571 // CHECK: call { i32, i32, i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n16.x4.b8x16.b4x16_p64.p3(ptr addrspace(3) %{{.*}})572 %m8n16_b4_l4 = nvvm.ldmatrix %arg0 {num = 4 : i32, layout = #nvvm.mma_layout<row>, shape = #nvvm.ld_st_matrix_shape<m = 8, n = 16>, eltType = #nvvm.ld_st_matrix_elt_type<b8x16.b4x16_p64>} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32, i32, i32)>573 574 // CHECK: call { i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m16n16.x1.trans.b8.p3(ptr addrspace(3) %{{.*}})575 %m16n16_l1t = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout<col>, shape =#nvvm.ld_st_matrix_shape<m = 16, n = 16>, eltType = #nvvm.ld_st_matrix_elt_type<b8>} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32)>576 // CHECK: call { i32, i32, i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m16n16.x2.trans.b8.p3(ptr addrspace(3) %{{.*}})577 %m16n16_l2t = nvvm.ldmatrix %arg0{num = 2 : i32, layout = #nvvm.mma_layout<col>, shape = #nvvm.ld_st_matrix_shape<m = 16, n = 16>,eltType =#nvvm.ld_st_matrix_elt_type<b8>} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32, i32, i32)>578 579 // CHECK: call { i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m16n16.x1.trans.b8x16.b6x16_p32.p3(ptr addrspace(3) %{{.*}})580 %m16n16_b6_l1t = nvvm.ldmatrix %arg0 {num = 1: i32, layout = #nvvm.mma_layout<col>, shape =#nvvm.ld_st_matrix_shape<m = 16, n = 16>, eltType = #nvvm.ld_st_matrix_elt_type<b8x16.b6x16_p32>} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32)>581 // CHECK: call { i32, i32, i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m16n16.x2.trans.b8x16.b6x16_p32.p3(ptr addrspace(3) %{{.*}})582 %m16n16_b6_l2t = nvvm.ldmatrix %arg0 {num = 2 : i32, layout = #nvvm.mma_layout<col>, shape =#nvvm.ld_st_matrix_shape<m = 16, n = 16>, eltType = #nvvm.ld_st_matrix_elt_type<b8x16.b6x16_p32>} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32, i32, i32)>583 584 // CHECK: call { i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m16n16.x1.trans.b8x16.b4x16_p64.p3(ptr addrspace(3) %{{.*}})585 %m16n16_b4_l1t = nvvm.ldmatrix %arg0 {num = 1: i32, layout = #nvvm.mma_layout<col>, shape =#nvvm.ld_st_matrix_shape<m = 16, n = 16>, eltType = #nvvm.ld_st_matrix_elt_type<b8x16.b4x16_p64>} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32)>586 // CHECK: call { i32, i32, i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m16n16.x2.trans.b8x16.b4x16_p64.p3(ptr addrspace(3) %{{.*}})587 %m16n16_b4_l2t = nvvm.ldmatrix %arg0 {num = 2 : i32, layout = #nvvm.mma_layout<col>, shape = #nvvm.ld_st_matrix_shape<m = 16, n = 16>,eltType =#nvvm.ld_st_matrix_elt_type<b8x16.b4x16_p64>} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32, i32, i32)>588 llvm.return589}590 591// CHECK-LABEL: @st_matrix592llvm.func @st_matrix(%arg0: !llvm.ptr<3>, %r1: i32, %r2: i32, %r3: i32, %r4: i32) {593 // CHECK: call void @llvm.nvvm.stmatrix.sync.aligned.m8n8.x1.b16.p3(ptr addrspace(3) %{{.*}}, i32 %{{.*}})594 nvvm.stmatrix %arg0, %r1 {layout = #nvvm.mma_layout<row>, shape = #nvvm.ld_st_matrix_shape<m = 8, n = 8>, eltType = #nvvm.ld_st_matrix_elt_type<b16>} : !llvm.ptr<3>, i32595 // CHECK: call void @llvm.nvvm.stmatrix.sync.aligned.m8n8.x1.trans.b16.p3(ptr addrspace(3) %{{.*}}, i32 %{{.*}})596 nvvm.stmatrix %arg0, %r1 {layout = #nvvm.mma_layout<col>, shape = #nvvm.ld_st_matrix_shape<m = 8, n = 8>, eltType = #nvvm.ld_st_matrix_elt_type<b16>} : !llvm.ptr<3>, i32597 // CHECK: call void @llvm.nvvm.stmatrix.sync.aligned.m16n8.x1.trans.b8.p3(ptr addrspace(3) %{{.*}}, i32 %{{.*}})598 nvvm.stmatrix %arg0, %r1 {layout = #nvvm.mma_layout<col>, shape = #nvvm.ld_st_matrix_shape<m = 16, n = 8>, eltType = #nvvm.ld_st_matrix_elt_type<b8>} : !llvm.ptr<3>, i32599 // CHECK: call void @llvm.nvvm.stmatrix.sync.aligned.m8n8.x2.b16.p3(ptr addrspace(3) %{{.*}}, i32 %{{.*}}, i32 %{{.*}})600 nvvm.stmatrix %arg0, %r1, %r2 {layout = #nvvm.mma_layout<row>, shape = #nvvm.ld_st_matrix_shape<m = 8, n = 8>, eltType = #nvvm.ld_st_matrix_elt_type<b16>} : !llvm.ptr<3>, i32, i32601 // CHECK: call void @llvm.nvvm.stmatrix.sync.aligned.m8n8.x2.trans.b16.p3(ptr addrspace(3) %{{.*}}, i32 %{{.*}}, i32 %{{.*}})602 nvvm.stmatrix %arg0, %r1, %r2 {layout = #nvvm.mma_layout<col>, shape = #nvvm.ld_st_matrix_shape<m = 8, n = 8>, eltType = #nvvm.ld_st_matrix_elt_type<b16>} : !llvm.ptr<3>, i32, i32603 // CHECK: call void @llvm.nvvm.stmatrix.sync.aligned.m16n8.x2.trans.b8.p3(ptr addrspace(3) %{{.*}}, i32 %{{.*}}, i32 %{{.*}})604 nvvm.stmatrix %arg0, %r1, %r2 {layout = #nvvm.mma_layout<col>, shape = #nvvm.ld_st_matrix_shape<m = 16, n = 8>, eltType = #nvvm.ld_st_matrix_elt_type<b8>} : !llvm.ptr<3>, i32, i32605 // CHECK: call void @llvm.nvvm.stmatrix.sync.aligned.m8n8.x4.b16.p3(ptr addrspace(3) %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}})606 nvvm.stmatrix %arg0, %r1, %r2, %r3, %r4 {layout = #nvvm.mma_layout<row>, shape = #nvvm.ld_st_matrix_shape<m = 8, n = 8>, eltType = #nvvm.ld_st_matrix_elt_type<b16>} : !llvm.ptr<3>, i32, i32, i32, i32607 // CHECK: call void @llvm.nvvm.stmatrix.sync.aligned.m8n8.x4.trans.b16.p3(ptr addrspace(3) %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}})608 nvvm.stmatrix %arg0, %r1, %r2, %r3, %r4 {layout = #nvvm.mma_layout<col>, shape = #nvvm.ld_st_matrix_shape<m = 8, n = 8>, eltType = #nvvm.ld_st_matrix_elt_type<b16>} : !llvm.ptr<3>, i32, i32, i32, i32609 // CHECK: call void @llvm.nvvm.stmatrix.sync.aligned.m16n8.x4.trans.b8.p3(ptr addrspace(3) %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 %{{.*}})610 nvvm.stmatrix %arg0, %r1, %r2, %r3, %r4 {layout = #nvvm.mma_layout<col>, shape = #nvvm.ld_st_matrix_shape<m = 16, n = 8>, eltType = #nvvm.ld_st_matrix_elt_type<b8>} : !llvm.ptr<3>, i32, i32, i32, i32611 llvm.return612}613 614// This function has the "kernel" attribute attached and should appear in the615// NVVM annotations after conversion.616llvm.func @kernel_func() attributes {nvvm.kernel} {617 llvm.return618}619 620// CHECK: ptx_kernel void @kernel_func621 622// -----623 624llvm.func @kernel_func() attributes {nvvm.kernel, nvvm.maxntid = array<i32: 1, 23, 32>} {625 llvm.return626}627 628// CHECK: define ptx_kernel void @kernel_func() #[[ATTR0:[0-9]+]]629// CHECK: attributes #[[ATTR0]] = { "nvvm.maxntid"="1,23,32" }630// -----631 632llvm.func @kernel_func() attributes {nvvm.kernel, nvvm.reqntid = array<i32: 1, 23, 32>} {633 llvm.return634}635 636// CHECK: define ptx_kernel void @kernel_func() #[[ATTR0:[0-9]+]]637// CHECK: attributes #[[ATTR0]] = { "nvvm.reqntid"="1,23,32" }638// -----639 640llvm.func @kernel_func() attributes {nvvm.kernel, nvvm.cluster_dim = array<i32: 3, 5, 7>} {641 llvm.return642}643 644// CHECK: define ptx_kernel void @kernel_func() #[[ATTR0:[0-9]+]]645// CHECK: attributes #[[ATTR0]] = { "nvvm.cluster_dim"="3,5,7" }646// -----647 648llvm.func @kernel_func() attributes {nvvm.kernel, nvvm.cluster_max_blocks = 8} {649 llvm.return650}651 652// CHECK: define ptx_kernel void @kernel_func() #[[ATTR0:[0-9]+]]653// CHECK: attributes #[[ATTR0]] = { "nvvm.maxclusterrank"="8" }654 655// -----656 657llvm.func @kernel_func() attributes {nvvm.kernel, nvvm.minctasm = 16} {658 llvm.return659}660 661// CHECK: define ptx_kernel void @kernel_func() #[[ATTR0:[0-9]+]]662// CHECK: attributes #[[ATTR0]] = { "nvvm.minctasm"="16" }663// -----664 665llvm.func @kernel_func() attributes {nvvm.kernel, nvvm.maxnreg = 16} {666 llvm.return667}668 669// CHECK: define ptx_kernel void @kernel_func() #[[ATTR0:[0-9]+]]670// CHECK: attributes #[[ATTR0]] = { "nvvm.maxnreg"="16" }671// -----672 673llvm.func @kernel_func() attributes {nvvm.kernel, nvvm.maxntid = array<i32: 1, 23, 32>,674 nvvm.minctasm = 16, nvvm.maxnreg = 32} {675 llvm.return676}677 678// CHECK: define ptx_kernel void @kernel_func() #[[ATTR0:[0-9]+]]679// CHECK: attributes #[[ATTR0]] = { "nvvm.maxnreg"="32" "nvvm.maxntid"="1,23,32" "nvvm.minctasm"="16" }680// -----681 682llvm.func @kernel_func() attributes {nvvm.kernel, nvvm.blocksareclusters,683 nvvm.reqntid = array<i32: 1, 23, 32>,684 nvvm.cluster_dim = array<i32: 3, 5, 7>} {685 llvm.return686}687 688// CHECK: define ptx_kernel void @kernel_func() #[[ATTR0:[0-9]+]]689// CHECK: attributes #[[ATTR0]] = { "nvvm.blocksareclusters" "nvvm.cluster_dim"="3,5,7" "nvvm.reqntid"="1,23,32" }690// -----691// CHECK: define ptx_kernel void @kernel_func(ptr byval(i32) "nvvm.grid_constant" %0)692llvm.func @kernel_func(%arg0: !llvm.ptr {llvm.byval = i32, nvvm.grid_constant}) attributes {nvvm.kernel} {693 llvm.return694}695 696// -----697// CHECK: define ptx_kernel void @kernel_func(ptr byval(i32) "nvvm.grid_constant" %0, float %1, ptr byval(float) "nvvm.grid_constant" %2)698llvm.func @kernel_func(%arg0: !llvm.ptr {llvm.byval = i32, nvvm.grid_constant}, %arg1: f32, %arg2: !llvm.ptr {llvm.byval = f32, nvvm.grid_constant}) attributes {nvvm.kernel} {699 llvm.return700}701 702 703// -----704// CHECK-LABEL: @nvvm_fence_proxy_tensormap_generic_release705llvm.func @nvvm_fence_proxy_tensormap_generic_release() {706 %c128 = llvm.mlir.constant(128) : i32707 // CHECK: call void @llvm.nvvm.fence.proxy.tensormap_generic.release.cta()708 nvvm.fence.proxy.release #nvvm.mem_scope<cta>709 710 // CHECK: call void @llvm.nvvm.fence.proxy.tensormap_generic.release.cluster()711 nvvm.fence.proxy.release #nvvm.mem_scope<cluster>712 713 // CHECK: call void @llvm.nvvm.fence.proxy.tensormap_generic.release.gpu()714 nvvm.fence.proxy.release #nvvm.mem_scope<gpu>715 716 // CHECK: call void @llvm.nvvm.fence.proxy.tensormap_generic.release.sys()717 nvvm.fence.proxy.release #nvvm.mem_scope<sys>718 llvm.return719}720 721// -----722// CHECK-LABEL: @nvvm_fence_proxy_tensormap_generic_acquire723llvm.func @nvvm_fence_proxy_tensormap_generic_acquire(%addr : !llvm.ptr) {724 %c128 = llvm.mlir.constant(128) : i32725 // CHECK: call void @llvm.nvvm.fence.proxy.tensormap_generic.acquire.cta(ptr {{%[0-9]+}}, i32 128)726 nvvm.fence.proxy.acquire #nvvm.mem_scope<cta> %addr, %c128727 728 // CHECK: call void @llvm.nvvm.fence.proxy.tensormap_generic.acquire.cluster(ptr {{%[0-9]+}}, i32 128)729 nvvm.fence.proxy.acquire #nvvm.mem_scope<cluster> %addr, %c128730 731 // CHECK: call void @llvm.nvvm.fence.proxy.tensormap_generic.acquire.gpu(ptr {{%[0-9]+}}, i32 128)732 nvvm.fence.proxy.acquire #nvvm.mem_scope<gpu> %addr, %c128733 734 // CHECK: call void @llvm.nvvm.fence.proxy.tensormap_generic.acquire.sys(ptr {{%[0-9]+}}, i32 128)735 nvvm.fence.proxy.acquire #nvvm.mem_scope<sys> %addr, %c128736 llvm.return737}738// -----739 740// CHECK-LABEL: @nvvm_exit741llvm.func @nvvm_exit() {742 // CHECK: call void @llvm.nvvm.exit()743 nvvm.exit744 llvm.return745}746 747 748 749// -----750// CHECK-LABEL: @nvvm_breakpoint751llvm.func @nvvm_breakpoint() {752 // CHECK: call void @llvm.debugtrap()753 nvvm.breakpoint754 llvm.return755}756 757// -----758// CHECK-LABEL: @nvvm_wgmma_fence_aligned759llvm.func @nvvm_wgmma_fence_aligned() {760 // CHECK: call void @llvm.nvvm.wgmma.fence.sync.aligned()761 nvvm.wgmma.fence.aligned762 llvm.return763}764 765// -----766// CHECK-LABEL: @nvvm_wgmma_commit_group_aligned767llvm.func @nvvm_wgmma_commit_group_aligned() {768 // CHECK: call void @llvm.nvvm.wgmma.commit_group.sync.aligned()769 nvvm.wgmma.commit.group.sync.aligned770 llvm.return771}772 773// -----774// CHECK-LABEL: @nvvm_wgmma_wait_group_aligned775llvm.func @nvvm_wgmma_wait_group_aligned() {776 // CHECK: call void @llvm.nvvm.wgmma.wait_group.sync.aligned(i64 0)777 nvvm.wgmma.wait.group.sync.aligned 0778 // CHECK: call void @llvm.nvvm.wgmma.wait_group.sync.aligned(i64 20)779 nvvm.wgmma.wait.group.sync.aligned 20780 llvm.return781}782 783// -----784// CHECK-LABEL: @nvvm_griddepcontrol_wait785llvm.func @nvvm_griddepcontrol_wait() {786 // CHECK: call void @llvm.nvvm.griddepcontrol.wait()787 nvvm.griddepcontrol wait788 llvm.return789}790 791// -----792// CHECK-LABEL: @nvvm_griddepcontrol_launch_dependents793llvm.func @nvvm_griddepcontrol_launch_dependents() {794 // CHECK: call void @llvm.nvvm.griddepcontrol.launch.dependents()795 nvvm.griddepcontrol launch_dependents796 llvm.return797}798 799// -----800// CHECK-LABEL: @nvvm_mapa801llvm.func @nvvm_mapa(%a: !llvm.ptr, %a_shared: !llvm.ptr<3>, %b : i32) {802 // CHECK-LLVM: call ptr @llvm.nvvm.mapa(ptr %{{.*}}, i32 %{{.*}})803 %0 = nvvm.mapa %a, %b: !llvm.ptr -> !llvm.ptr804 // CHECK-LLVM: call ptr addrspace(7) @llvm.nvvm.mapa.shared.cluster(ptr addrspace(3) %{{.*}}, i32 %{{.*}})805 %1 = nvvm.mapa %a_shared, %b: !llvm.ptr<3> -> !llvm.ptr<7>806 llvm.return807}808 809// -----810// CHECK-LABEL: @nvvm_redux_sync811llvm.func @nvvm_redux_sync(%value: i32, %offset: i32) {812 // CHECK: call i32 @llvm.nvvm.redux.sync.add(i32 %{{.*}}, i32 %{{.*}})813 %0 = nvvm.redux.sync add %value, %offset: i32 -> i32814 // CHECK: call i32 @llvm.nvvm.redux.sync.umax(i32 %{{.*}}, i32 %{{.*}})815 %1 = nvvm.redux.sync umax %value, %offset: i32 -> i32816 // CHECK: call i32 @llvm.nvvm.redux.sync.umin(i32 %{{.*}}, i32 %{{.*}})817 %2 = nvvm.redux.sync umin %value, %offset: i32 -> i32818 // CHECK: call i32 @llvm.nvvm.redux.sync.and(i32 %{{.*}}, i32 %{{.*}})819 %3 = nvvm.redux.sync and %value, %offset: i32 -> i32820 // CHECK: call i32 @llvm.nvvm.redux.sync.or(i32 %{{.*}}, i32 %{{.*}})821 %4 = nvvm.redux.sync or %value, %offset: i32 -> i32822 // CHECK: call i32 @llvm.nvvm.redux.sync.xor(i32 %{{.*}}, i32 %{{.*}})823 %5 = nvvm.redux.sync xor %value, %offset: i32 -> i32824 // CHECK: call i32 @llvm.nvvm.redux.sync.max(i32 %{{.*}}, i32 %{{.*}})825 %6 = nvvm.redux.sync max %value, %offset: i32 -> i32826 // CHECK: call i32 @llvm.nvvm.redux.sync.min(i32 %{{.*}}, i32 %{{.*}})827 %7 = nvvm.redux.sync min %value, %offset: i32 -> i32828 llvm.return829}830 831// CHECK-LABEL: @nvvm_redux_sync_f32832llvm.func @nvvm_redux_sync_f32(%value: f32, %offset: i32) {833 // CHECK: call float @llvm.nvvm.redux.sync.fmin(float %{{.*}}, i32 %{{.*}})834 %0 = nvvm.redux.sync fmin %value, %offset: f32 -> f32835 // CHECK: call float @llvm.nvvm.redux.sync.fmin.abs(float %{{.*}}, i32 %{{.*}})836 %1 = nvvm.redux.sync fmin %value, %offset {abs = true}: f32 -> f32837 // CHECK: call float @llvm.nvvm.redux.sync.fmin.NaN(float %{{.*}}, i32 %{{.*}})838 %2 = nvvm.redux.sync fmin %value, %offset {nan = true}: f32 -> f32839 // CHECK: call float @llvm.nvvm.redux.sync.fmin.abs.NaN(float %{{.*}}, i32 %{{.*}})840 %3 = nvvm.redux.sync fmin %value, %offset {abs = true, nan = true}: f32 -> f32841 // CHECK: call float @llvm.nvvm.redux.sync.fmax(float %{{.*}}, i32 %{{.*}})842 %4 = nvvm.redux.sync fmax %value, %offset: f32 -> f32843 // CHECK: call float @llvm.nvvm.redux.sync.fmax.abs(float %{{.*}}, i32 %{{.*}})844 %5 = nvvm.redux.sync fmax %value, %offset {abs = true}: f32 -> f32845 // CHECK: call float @llvm.nvvm.redux.sync.fmax.NaN(float %{{.*}}, i32 %{{.*}})846 %6 = nvvm.redux.sync fmax %value, %offset {nan = true}: f32 -> f32847 // CHECK: call float @llvm.nvvm.redux.sync.fmax.abs.NaN(float %{{.*}}, i32 %{{.*}})848 %7 = nvvm.redux.sync fmax %value, %offset {abs = true, nan = true}: f32 -> f32849 llvm.return850}851 852// -----853// CHECK-LABEL: @nvvm_match_sync854llvm.func @nvvm_match_sync(%mask: i32, %val32: i32, %val64: i64) {855 // CHECK: call i32 @llvm.nvvm.match.any.sync.i32(i32 %{{.*}}, i32 %{{.*}})856 %0 = nvvm.match.sync any %mask, %val32 : i32 -> i32857 // CHECK: call { i32, i1 } @llvm.nvvm.match.all.sync.i32p(i32 %{{.*}}, i32 %{{.*}})858 %1 = nvvm.match.sync all %mask, %val32 : i32 -> !llvm.struct<(i32, i1)>859 // CHECK: call i32 @llvm.nvvm.match.any.sync.i64(i32 %{{.*}}, i64 %{{.*}})860 %2 = nvvm.match.sync any %mask, %val64 : i64 -> i32861 // CHECK: call { i32, i1 } @llvm.nvvm.match.all.sync.i64p(i32 %{{.*}}, i64 %{{.*}})862 %3 = nvvm.match.sync all %mask, %val64 : i64 -> !llvm.struct<(i32, i1)>863 llvm.return864}865 866// -----867// CHECK-LABEL: @nvvm_st_bulk868llvm.func @nvvm_st_bulk(%addr_gen: !llvm.ptr, %addr_shared: !llvm.ptr<3>, %size: i64) {869 // CHECK: call void @llvm.nvvm.st.bulk(ptr %{{.*}}, i64 %{{.*}}, i64 0)870 nvvm.st.bulk %addr_gen, size = %size : !llvm.ptr871 // CHECK: call void @llvm.nvvm.st.bulk.shared.cta(ptr addrspace(3) %{{.*}}, i64 %{{.*}}, i64 0)872 nvvm.st.bulk %addr_shared, size = %size: !llvm.ptr<3>873 // CHECK: call void @llvm.nvvm.st.bulk(ptr %{{.*}}, i64 %{{.*}}, i64 0)874 nvvm.st.bulk %addr_gen, size = %size, init = 0 : !llvm.ptr875 // CHECK: call void @llvm.nvvm.st.bulk.shared.cta(ptr addrspace(3) %{{.*}}, i64 %{{.*}}, i64 0)876 nvvm.st.bulk %addr_shared, size = %size, init = 0: !llvm.ptr<3>877 llvm.return878}879 880// -----881// CHECK-LABEL: @nvvm_dot_accumulate_4way882llvm.func @nvvm_dot_accumulate_4way(%a: vector<4xi8>, %b: vector<4xi8>, %c: i32) {883 // CHECK: %[[a_cast:.*]] = bitcast <4 x i8> %{{.*}} to i32884 // CHECK: %[[b_cast:.*]] = bitcast <4 x i8> %{{.*}} to i32885 // CHECK: call i32 @llvm.nvvm.idp4a.u.u(i32 %[[a_cast]], i32 %[[b_cast]], i32 %{{.*}})886 %0 = nvvm.dot.accumulate.4way %a <unsigned>, %b <unsigned>, %c: vector<4xi8>, vector<4xi8>887 // CHECK: %[[a_cast:.*]] = bitcast <4 x i8> %{{.*}} to i32888 // CHECK: %[[b_cast:.*]] = bitcast <4 x i8> %{{.*}} to i32889 // CHECK: call i32 @llvm.nvvm.idp4a.s.u(i32 %[[a_cast]], i32 %[[b_cast]], i32 %{{.*}})890 %1 = nvvm.dot.accumulate.4way %a <signed>, %b <unsigned>, %c: vector<4xi8>, vector<4xi8>891 // CHECK: %[[a_cast:.*]] = bitcast <4 x i8> %{{.*}} to i32892 // CHECK: %[[b_cast:.*]] = bitcast <4 x i8> %{{.*}} to i32893 // CHECK: call i32 @llvm.nvvm.idp4a.u.s(i32 %[[a_cast]], i32 %[[b_cast]], i32 %{{.*}})894 %2 = nvvm.dot.accumulate.4way %a <unsigned>, %b <signed>, %c: vector<4xi8>, vector<4xi8>895 // CHECK: %[[a_cast:.*]] = bitcast <4 x i8> %{{.*}} to i32896 // CHECK: %[[b_cast:.*]] = bitcast <4 x i8> %{{.*}} to i32897 // CHECK: call i32 @llvm.nvvm.idp4a.s.s(i32 %[[a_cast]], i32 %[[b_cast]], i32 %{{.*}})898 %3 = nvvm.dot.accumulate.4way %a <signed>, %b <signed>, %c: vector<4xi8>, vector<4xi8>899 llvm.return900}901 902// -----903// CHECK-LABEL: @nvvm_dot_accumulate_2way904llvm.func @nvvm_dot_accumulate_2way(%a: vector<2xi16>, %b: vector<4xi8>, %c: i32) {905 // CHECK: %[[a_cast:.*]] = bitcast <2 x i16> %{{.*}} to i32906 // CHECK: %[[b_cast:.*]] = bitcast <4 x i8> %{{.*}} to i32907 // CHECK: call i32 @llvm.nvvm.idp2a.u.u(i32 %[[a_cast]], i32 %[[b_cast]], i1 false, i32 %{{.*}})908 %0 = nvvm.dot.accumulate.2way %a <unsigned>, %b <unsigned>, %c {b_hi = false} : vector<2xi16>, vector<4xi8>909 // CHECK: %[[a_cast:.*]] = bitcast <2 x i16> %{{.*}} to i32910 // CHECK: %[[b_cast:.*]] = bitcast <4 x i8> %{{.*}} to i32911 // CHECK: call i32 @llvm.nvvm.idp2a.u.u(i32 %[[a_cast]], i32 %[[b_cast]], i1 true, i32 %{{.*}})912 %1 = nvvm.dot.accumulate.2way %a <unsigned>, %b <unsigned>, %c {b_hi = true}: vector<2xi16>, vector<4xi8>913 // CHECK: %[[a_cast:.*]] = bitcast <2 x i16> %{{.*}} to i32914 // CHECK: %[[b_cast:.*]] = bitcast <4 x i8> %{{.*}} to i32915 // CHECK: call i32 @llvm.nvvm.idp2a.s.u(i32 %[[a_cast]], i32 %[[b_cast]], i1 false, i32 %{{.*}})916 %2 = nvvm.dot.accumulate.2way %a <signed>, %b <unsigned>, %c {b_hi = false}: vector<2xi16>, vector<4xi8>917 // CHECK: %[[a_cast:.*]] = bitcast <2 x i16> %{{.*}} to i32918 // CHECK: %[[b_cast:.*]] = bitcast <4 x i8> %{{.*}} to i32919 // CHECK: call i32 @llvm.nvvm.idp2a.s.u(i32 %[[a_cast]], i32 %[[b_cast]], i1 true, i32 %{{.*}})920 %3 = nvvm.dot.accumulate.2way %a <signed>, %b <unsigned>, %c {b_hi = true}: vector<2xi16>, vector<4xi8>921 // CHECK: %[[a_cast:.*]] = bitcast <2 x i16> %{{.*}} to i32922 // CHECK: %[[b_cast:.*]] = bitcast <4 x i8> %{{.*}} to i32923 // CHECK: call i32 @llvm.nvvm.idp2a.u.s(i32 %[[a_cast]], i32 %[[b_cast]], i1 false, i32 %{{.*}})924 %4 = nvvm.dot.accumulate.2way %a <unsigned>, %b <signed>, %c {b_hi = false}: vector<2xi16>, vector<4xi8>925 // CHECK: %[[a_cast:.*]] = bitcast <2 x i16> %{{.*}} to i32926 // CHECK: %[[b_cast:.*]] = bitcast <4 x i8> %{{.*}} to i32927 // CHECK: call i32 @llvm.nvvm.idp2a.u.s(i32 %[[a_cast]], i32 %[[b_cast]], i1 true, i32 %{{.*}})928 %5 = nvvm.dot.accumulate.2way %a <unsigned>, %b <signed>, %c {b_hi = true}: vector<2xi16>, vector<4xi8>929 // CHECK: %[[a_cast:.*]] = bitcast <2 x i16> %{{.*}} to i32930 // CHECK: %[[b_cast:.*]] = bitcast <4 x i8> %{{.*}} to i32931 // CHECK: call i32 @llvm.nvvm.idp2a.s.s(i32 %[[a_cast]], i32 %[[b_cast]], i1 false, i32 %{{.*}})932 %6 = nvvm.dot.accumulate.2way %a <signed>, %b <signed>, %c {b_hi = false}: vector<2xi16>, vector<4xi8>933 // CHECK: %[[a_cast:.*]] = bitcast <2 x i16> %{{.*}} to i32934 // CHECK: %[[b_cast:.*]] = bitcast <4 x i8> %{{.*}} to i32935 // CHECK: call i32 @llvm.nvvm.idp2a.s.s(i32 %[[a_cast]], i32 %[[b_cast]], i1 true, i32 %{{.*}})936 %7 = nvvm.dot.accumulate.2way %a <signed>, %b <signed>, %c {b_hi = true}: vector<2xi16>, vector<4xi8>937 llvm.return938}939 940// -----941 942// CHECK-LABEL: @nvvm_pmevent943llvm.func @nvvm_pmevent() {944 // CHECK: call void @llvm.nvvm.pm.event.mask(i16 15000)945 nvvm.pmevent mask = 15000946 // CHECK: call void @llvm.nvvm.pm.event.mask(i16 4)947 nvvm.pmevent mask = 4948 llvm.return949}950 951// -----952 953// CHECK-LABEL: @nanosleep954llvm.func @nanosleep(%duration: i32) {955 // CHECK: call void @llvm.nvvm.nanosleep(i32 %{{.*}})956 nvvm.nanosleep %duration957 llvm.return958}959