361 lines · plain
1; RUN: mlir-translate -import-llvm %s | FileCheck %s2 3; CHECK-LABEL: @nvvm_special_regs4define i32 @nvvm_special_regs() {5 ; CHECK: = nvvm.read.ptx.sreg.tid.x : i326 %1 = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()7 ; CHECK: = nvvm.read.ptx.sreg.tid.y : i328 %2 = call i32 @llvm.nvvm.read.ptx.sreg.tid.y()9 ; CHECK: = nvvm.read.ptx.sreg.tid.z : i3210 %3 = call i32 @llvm.nvvm.read.ptx.sreg.tid.z()11 ; CHECK: = nvvm.read.ptx.sreg.ntid.x : i3212 %4 = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()13 ; CHECK: = nvvm.read.ptx.sreg.ntid.y : i3214 %5 = call i32 @llvm.nvvm.read.ptx.sreg.ntid.y()15 ; CHECK: = nvvm.read.ptx.sreg.ntid.z : i3216 %6 = call i32 @llvm.nvvm.read.ptx.sreg.ntid.z()17 ; CHECK: = nvvm.read.ptx.sreg.ctaid.x : i3218 %7 = call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()19 ; CHECK: = nvvm.read.ptx.sreg.ctaid.y : i3220 %8 = call i32 @llvm.nvvm.read.ptx.sreg.ctaid.y()21 ; CHECK: = nvvm.read.ptx.sreg.ctaid.z : i3222 %9 = call i32 @llvm.nvvm.read.ptx.sreg.ctaid.z()23 ; CHECK: = nvvm.read.ptx.sreg.nctaid.x : i3224 %10 = call i32 @llvm.nvvm.read.ptx.sreg.nctaid.x()25 ; CHECK: = nvvm.read.ptx.sreg.nctaid.y : i3226 %11 = call i32 @llvm.nvvm.read.ptx.sreg.nctaid.y()27 ; CHECK: = nvvm.read.ptx.sreg.nctaid.z : i3228 %12 = call i32 @llvm.nvvm.read.ptx.sreg.nctaid.z()29 ; CHECK: = nvvm.read.ptx.sreg.warpsize : i3230 %13 = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()31 ; CHECK: = nvvm.read.ptx.sreg.laneid : i3232 %14 = call i32 @llvm.nvvm.read.ptx.sreg.laneid()33 ; CHECK: = nvvm.read.ptx.sreg.clusterid.x : i3234 %15 = call i32 @llvm.nvvm.read.ptx.sreg.clusterid.x()35 ; CHECK: = nvvm.read.ptx.sreg.clusterid.y : i3236 %16 = call i32 @llvm.nvvm.read.ptx.sreg.clusterid.y()37 ; CHECK: = nvvm.read.ptx.sreg.clusterid.z : i3238 %17 = call i32 @llvm.nvvm.read.ptx.sreg.clusterid.z()39 ; CHECK: = nvvm.read.ptx.sreg.nclusterid.x : i3240 %18 = call i32 @llvm.nvvm.read.ptx.sreg.nclusterid.x()41 ; CHECK: = nvvm.read.ptx.sreg.nclusterid.y : i3242 %19 = call i32 @llvm.nvvm.read.ptx.sreg.nclusterid.y()43 ; CHECK: = nvvm.read.ptx.sreg.nclusterid.z : i3244 %20 = call i32 @llvm.nvvm.read.ptx.sreg.nclusterid.z()45 ; CHECK: = nvvm.read.ptx.sreg.cluster.ctaid.x : i3246 %21 = call i32 @llvm.nvvm.read.ptx.sreg.cluster.ctaid.x()47 ; CHECK: = nvvm.read.ptx.sreg.cluster.ctaid.y : i3248 %22 = call i32 @llvm.nvvm.read.ptx.sreg.cluster.ctaid.y()49 ; CHECK: = nvvm.read.ptx.sreg.cluster.ctaid.z : i3250 %23 = call i32 @llvm.nvvm.read.ptx.sreg.cluster.ctaid.z()51 ; CHECK: = nvvm.read.ptx.sreg.cluster.nctaid.x : i3252 %24 = call i32 @llvm.nvvm.read.ptx.sreg.cluster.nctaid.x()53 ; CHECK: = nvvm.read.ptx.sreg.cluster.nctaid.y : i3254 %25 = call i32 @llvm.nvvm.read.ptx.sreg.cluster.nctaid.y()55 ; CHECK: = nvvm.read.ptx.sreg.cluster.nctaid.z : i3256 %26 = call i32 @llvm.nvvm.read.ptx.sreg.cluster.nctaid.z()57 ; CHECK: = nvvm.read.ptx.sreg.cluster.ctarank : i3258 %27 = call i32 @llvm.nvvm.read.ptx.sreg.cluster.ctarank()59 ; CHECK: = nvvm.read.ptx.sreg.cluster.nctarank : i3260 %28 = call i32 @llvm.nvvm.read.ptx.sreg.cluster.nctarank()61 62 ; CHECK = nvvm.read.ptx.sreg.tid.x range <0 : i32, 64 : i32> : i3263 %29 = call range(i32 0, 64) i32 @llvm.nvvm.read.ptx.sreg.tid.x()64 ret i32 %165}66 67; CHECK-LABEL: @nvvm_rcp68define float @nvvm_rcp(float %0) {69 ; CHECK: = nvvm.rcp.approx.ftz.f %{{.*}} : f3270 %2 = call float @llvm.nvvm.rcp.approx.ftz.f(float %0)71 ret float %272}73 74; CHECK-LABEL: @llvm_nvvm_barrier0()75define void @llvm_nvvm_barrier0() {76 ; CHECK: llvm.nvvm.barrier.cta.sync.aligned.all77 call void @llvm.nvvm.barrier0()78 ret void79}80 81; TODO: Support the intrinsics below once they derive from NVVM_IntrOp rather than from NVVM_Op.82;83; define i32 @nvvm_shfl(i32 %0, i32 %1, i32 %2, i32 %3, float %4) {84; %6 = call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 %0, i32 %3, i32 %1, i32 %2)85; %7 = call float @llvm.nvvm.shfl.sync.bfly.f32(i32 %0, float %4, i32 %1, i32 %2)86; %8 = call i32 @llvm.nvvm.shfl.sync.up.i32(i32 %0, i32 %3, i32 %1, i32 %2)87; %9 = call float @llvm.nvvm.shfl.sync.up.f32(i32 %0, float %4, i32 %1, i32 %2)88; %10 = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 %0, i32 %3, i32 %1, i32 %2)89; %11 = call float @llvm.nvvm.shfl.sync.down.f32(i32 %0, float %4, i32 %1, i32 %2)90; %12 = call i32 @llvm.nvvm.shfl.sync.idx.i32(i32 %0, i32 %3, i32 %1, i32 %2)91; %13 = call float @llvm.nvvm.shfl.sync.idx.f32(i32 %0, float %4, i32 %1, i32 %2)92; ret i32 %693; }94;95; define { i32, i1 } @nvvm_shfl_pred(i32 %0, i32 %1, i32 %2, i32 %3, float %4) {96; %6 = call { i32, i1 } @llvm.nvvm.shfl.sync.bfly.i32p(i32 %0, i32 %3, i32 %1, i32 %2)97; %7 = call { float, i1 } @llvm.nvvm.shfl.sync.bfly.f32p(i32 %0, float %4, i32 %1, i32 %2)98; %8 = call { i32, i1 } @llvm.nvvm.shfl.sync.up.i32p(i32 %0, i32 %3, i32 %1, i32 %2)99; %9 = call { float, i1 } @llvm.nvvm.shfl.sync.up.f32p(i32 %0, float %4, i32 %1, i32 %2)100; %10 = call { i32, i1 } @llvm.nvvm.shfl.sync.down.i32p(i32 %0, i32 %3, i32 %1, i32 %2)101; %11 = call { float, i1 } @llvm.nvvm.shfl.sync.down.f32p(i32 %0, float %4, i32 %1, i32 %2)102; %12 = call { i32, i1 } @llvm.nvvm.shfl.sync.idx.i32p(i32 %0, i32 %3, i32 %1, i32 %2)103; %13 = call { float, i1 } @llvm.nvvm.shfl.sync.idx.f32p(i32 %0, float %4, i32 %1, i32 %2)104; ret { i32, i1 } %6105; }106;107; define i32 @nvvm_vote(i32 %0, i1 %1) {108; %3 = call i32 @llvm.nvvm.vote.ballot.sync(i32 %0, i1 %1)109; ret i32 %3110; }111;112; define { float, float, float, float, float, float, float, float } @nvvm_mma_mn8n8k4_row_col_f32_f32(<2 x half> %0, <2 x half> %1, <2 x half> %2, <2 x half> %3, float %4, float %5, float %6, float %7, float %8, float %9, float %10, float %11) {113; %13 = call { float, float, float, float, float, float, float, float } @llvm.nvvm.mma.m8n8k4.row.col.f32.f32(<2 x half> %0, <2 x half> %1, <2 x half> %2, <2 x half> %3, float %4, float %5, float %6, float %7, float %8, float %9, float %10, float %11)114; ret { float, float, float, float, float, float, float, float } %13115; }116;117; define { <2 x half>, <2 x half> } @nvvm_mma_m16n8k16_f16_f16(<2 x half> %0, <2 x half> %1, <2 x half> %2, <2 x half> %3, <2 x half> %4, <2 x half> %5, <2 x half> %6, <2 x half> %7) {118; %9 = call { <2 x half>, <2 x half> } @llvm.nvvm.mma.m16n8k16.row.col.f16.f16(<2 x half> %0, <2 x half> %1, <2 x half> %2, <2 x half> %3, <2 x half> %4, <2 x half> %5, <2 x half> %6, <2 x half> %7)119; ret { <2 x half>, <2 x half> } %9120; }121;122; define { float, float, float, float } @nvvm_mma_m16n8k16_f32_f16(<2 x half> %0, <2 x half> %1, <2 x half> %2, <2 x half> %3, <2 x half> %4, <2 x half> %5, <2 x half> %6, <2 x half> %7) {123; %9 = call { float, float, float, float } @llvm.nvvm.mma.m16n8k16.row.col.f32.f16(<2 x half> %0, <2 x half> %1, <2 x half> %2, <2 x half> %3, <2 x half> %4, <2 x half> %5, <2 x half> %6, <2 x half> %7)124; ret { float, float, float, float } %9125; }126;127; define { <2 x half>, <2 x half> } @nvvm_mma_m16n8k16_f16_f32(<2 x half> %0, <2 x half> %1, <2 x half> %2, <2 x half> %3, <2 x half> %4, <2 x half> %5, float %6, float %7, float %8, float %9) {128; %11 = call { <2 x half>, <2 x half> } @llvm.nvvm.mma.m16n8k16.row.col.f16.f32(<2 x half> %0, <2 x half> %1, <2 x half> %2, <2 x half> %3, <2 x half> %4, <2 x half> %5, float %6, float %7, float %8, float %9)129; ret { <2 x half>, <2 x half> } %11130; }131;132; define { float, float, float, float } @nvvm_mma_m16n8k16_f32_f32(<2 x half> %0, <2 x half> %1, <2 x half> %2, <2 x half> %3, <2 x half> %4, <2 x half> %5, float %6, float %7, float %8, float %9) {133; %11 = call { float, float, float, float } @llvm.nvvm.mma.m16n8k16.row.col.f32.f32(<2 x half> %0, <2 x half> %1, <2 x half> %2, <2 x half> %3, <2 x half> %4, <2 x half> %5, float %6, float %7, float %8, float %9)134; ret { float, float, float, float } %11135; }136;137; define { i32, i32, i32, i32 } @nvvm_mma_m16n8k16_s8_s8(i32 %0, i32 %1, i32 %2, i32 %3, i32 %4, i32 %5, i32 %6) {138; %8 = call { i32, i32, i32, i32 } @llvm.nvvm.mma.m16n8k16.row.col.s8(i32 %0, i32 %1, i32 %2, i32 %3, i32 %4, i32 %5, i32 %6)139; ret { i32, i32, i32, i32 } %8140; }141;142; define { i32, i32, i32, i32 } @nvvm_mma_m16n8k16_s8_u8(i32 %0, i32 %1, i32 %2, i32 %3, i32 %4, i32 %5, i32 %6) {143; %8 = call { i32, i32, i32, i32 } @llvm.nvvm.mma.m16n8k16.row.col.satfinite.s8.u8(i32 %0, i32 %1, i32 %2, i32 %3, i32 %4, i32 %5, i32 %6)144; ret { i32, i32, i32, i32 } %8145; }146;147; define { i32, i32, i32, i32 } @nvvm_mma_m16n8k128_b1_b1(i32 %0, i32 %1, i32 %2, i32 %3, i32 %4, i32 %5, i32 %6) {148; %8 = call { i32, i32, i32, i32 } @llvm.nvvm.mma.xor.popc.m16n8k128.row.col.b1(i32 %0, i32 %1, i32 %2, i32 %3, i32 %4, i32 %5, i32 %6)149; ret { i32, i32, i32, i32 } %8150; }151;152; define { i32, i32, i32, i32 } @nvvm_mma_m16n8k32_s4_s4(i32 %0, i32 %1, i32 %2, i32 %3, i32 %4, i32 %5, i32 %6) {153; %8 = call { i32, i32, i32, i32 } @llvm.nvvm.mma.m16n8k32.row.col.satfinite.s4(i32 %0, i32 %1, i32 %2, i32 %3, i32 %4, i32 %5, i32 %6)154; ret { i32, i32, i32, i32 } %8155; }156;157; define { double, double } @nvvm_mma_m8n8k4_f64_f64(double %0, double %1, double %2, double %3) {158; %5 = call { double, double } @llvm.nvvm.mma.m8n8k4.row.col.f64(double %0, double %1, double %2, double %3)159; ret { double, double } %5160; }161;162; define { float, float, float, float } @nvvm_mma_m16n8k4_tf32_f32(i32 %0, i32 %1, i32 %2, float %3, float %4, float %5, float %6) {163; %8 = call { float, float, float, float } @llvm.nvvm.mma.m16n8k4.row.col.tf32(i32 %0, i32 %1, i32 %2, float %3, float %4, float %5, float %6)164; ret { float, float, float, float } %8165; }166;167; define void @gpu_wmma_load_op(ptr addrspace(3) %0, i32 %1) {168; %3 = 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) %0, i32 %1)169; ret void170; }171;172; define void @gpu_wmma_store_op(ptr addrspace(3) %0, i32 %1, <2 x half> %2, <2 x half> %3, <2 x half> %4, <2 x half> %5) {173; call void @llvm.nvvm.wmma.m16n16k16.store.d.row.stride.f16.p3(ptr addrspace(3) %0, <2 x half> %2, <2 x half> %3, <2 x half> %4, <2 x half> %5, i32 %1)174; ret void175; }176;177; define void @gpu_wmma_mma_op(<2 x half> %0, <2 x half> %1, <2 x half> %2, <2 x half> %3, <2 x half> %4, <2 x half> %5, <2 x half> %6, <2 x half> %7, <2 x half> %8, <2 x half> %9, <2 x half> %10, <2 x half> %11, <2 x half> %12, <2 x half> %13, <2 x half> %14, <2 x half> %15, <2 x half> %16, <2 x half> %17, <2 x half> %18, <2 x half> %19) {178; %21 = 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> %0, <2 x half> %1, <2 x half> %2, <2 x half> %3, <2 x half> %4, <2 x half> %5, <2 x half> %6, <2 x half> %7, <2 x half> %8, <2 x half> %9, <2 x half> %10, <2 x half> %11, <2 x half> %12, <2 x half> %13, <2 x half> %14, <2 x half> %15, <2 x half> %16, <2 x half> %17, <2 x half> %18, <2 x half> %19)179; ret void180; }181;182; define void @nvvm_wmma_load_tf32(ptr %0, i32 %1) {183; %3 = call { i32, i32, i32, i32 } @llvm.nvvm.wmma.m16n16k8.load.a.row.stride.tf32.p0(ptr %0, i32 %1)184; ret void185; }186;187; define void @nvvm_wmma_mma(i32 %0, i32 %1, i32 %2, i32 %3, i32 %4, i32 %5, i32 %6, i32 %7, float %8, float %9, float %10, float %11, float %12, float %13, float %14, float %15) {188; %17 = call { float, float, float, float, float, float, float, float } @llvm.nvvm.wmma.m16n16k8.mma.row.row.tf32(i32 %0, i32 %1, i32 %2, i32 %3, i32 %4, i32 %5, i32 %6, i32 %7, float %8, float %9, float %10, float %11, float %12, float %13, float %14, float %15)189; ret void190; }191;192; define void @cp_async(ptr addrspace(3) %0, ptr addrspace(1) %1) {193; call void @llvm.nvvm.cp.async.ca.shared.global.4(ptr addrspace(3) %0, ptr addrspace(1) %1)194; call void @llvm.nvvm.cp.async.ca.shared.global.8(ptr addrspace(3) %0, ptr addrspace(1) %1)195; call void @llvm.nvvm.cp.async.ca.shared.global.16(ptr addrspace(3) %0, ptr addrspace(1) %1)196; call void @llvm.nvvm.cp.async.cg.shared.global.16(ptr addrspace(3) %0, ptr addrspace(1) %1)197; call void @llvm.nvvm.cp.async.commit.group()198; call void @llvm.nvvm.cp.async.wait.group(i32 0)199; ret void200; }201;202; define void @ld_matrix(ptr addrspace(3) %0) {203; %2 = call i32 @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x1.b16.p3(ptr addrspace(3) %0)204; %3 = call { i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x2.b16.p3(ptr addrspace(3) %0)205; %4 = call { i32, i32, i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x4.b16.p3(ptr addrspace(3) %0)206; %5 = call i32 @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x1.trans.b16.p3(ptr addrspace(3) %0)207; %6 = call { i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x2.trans.b16.p3(ptr addrspace(3) %0)208; %7 = call { i32, i32, i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x4.trans.b16.p3(ptr addrspace(3) %0)209; ret void210; }211 212declare noundef i32 @llvm.nvvm.read.ptx.sreg.tid.x()213 214declare noundef i32 @llvm.nvvm.read.ptx.sreg.tid.y()215 216declare noundef i32 @llvm.nvvm.read.ptx.sreg.tid.z()217 218declare noundef i32 @llvm.nvvm.read.ptx.sreg.ntid.x()219 220declare noundef i32 @llvm.nvvm.read.ptx.sreg.ntid.y()221 222declare noundef i32 @llvm.nvvm.read.ptx.sreg.ntid.z()223 224declare noundef i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()225 226declare noundef i32 @llvm.nvvm.read.ptx.sreg.ctaid.y()227 228declare noundef i32 @llvm.nvvm.read.ptx.sreg.ctaid.z()229 230declare noundef i32 @llvm.nvvm.read.ptx.sreg.nctaid.x()231 232declare noundef i32 @llvm.nvvm.read.ptx.sreg.nctaid.y()233 234declare noundef i32 @llvm.nvvm.read.ptx.sreg.nctaid.z()235 236declare noundef i32 @llvm.nvvm.read.ptx.sreg.warpsize()237 238declare noundef i32 @llvm.nvvm.read.ptx.sreg.laneid()239 240declare noundef i32 @llvm.nvvm.read.ptx.sreg.clusterid.x()241 242declare noundef i32 @llvm.nvvm.read.ptx.sreg.clusterid.y()243 244declare noundef i32 @llvm.nvvm.read.ptx.sreg.clusterid.z()245 246declare noundef i32 @llvm.nvvm.read.ptx.sreg.nclusterid.x()247 248declare noundef i32 @llvm.nvvm.read.ptx.sreg.nclusterid.y()249 250declare noundef i32 @llvm.nvvm.read.ptx.sreg.nclusterid.z()251 252declare noundef i32 @llvm.nvvm.read.ptx.sreg.cluster.ctaid.x()253 254declare noundef i32 @llvm.nvvm.read.ptx.sreg.cluster.ctaid.y()255 256declare noundef i32 @llvm.nvvm.read.ptx.sreg.cluster.ctaid.z()257 258declare noundef i32 @llvm.nvvm.read.ptx.sreg.cluster.nctaid.x()259 260declare noundef i32 @llvm.nvvm.read.ptx.sreg.cluster.nctaid.y()261 262declare noundef i32 @llvm.nvvm.read.ptx.sreg.cluster.nctaid.z()263 264declare noundef i32 @llvm.nvvm.read.ptx.sreg.cluster.ctarank()265 266declare noundef i32 @llvm.nvvm.read.ptx.sreg.cluster.nctarank()267 268declare float @llvm.nvvm.rcp.approx.ftz.f(float)269 270declare void @llvm.nvvm.barrier0()271 272declare i32 @llvm.nvvm.shfl.sync.bfly.i32(i32, i32, i32, i32)273 274declare float @llvm.nvvm.shfl.sync.bfly.f32(i32, float, i32, i32)275 276declare i32 @llvm.nvvm.shfl.sync.up.i32(i32, i32, i32, i32)277 278declare float @llvm.nvvm.shfl.sync.up.f32(i32, float, i32, i32)279 280declare i32 @llvm.nvvm.shfl.sync.down.i32(i32, i32, i32, i32)281 282declare float @llvm.nvvm.shfl.sync.down.f32(i32, float, i32, i32)283 284declare i32 @llvm.nvvm.shfl.sync.idx.i32(i32, i32, i32, i32)285 286declare float @llvm.nvvm.shfl.sync.idx.f32(i32, float, i32, i32)287 288declare { i32, i1 } @llvm.nvvm.shfl.sync.bfly.i32p(i32, i32, i32, i32)289 290declare { float, i1 } @llvm.nvvm.shfl.sync.bfly.f32p(i32, float, i32, i32)291 292declare { i32, i1 } @llvm.nvvm.shfl.sync.up.i32p(i32, i32, i32, i32)293 294declare { float, i1 } @llvm.nvvm.shfl.sync.up.f32p(i32, float, i32, i32)295 296declare { i32, i1 } @llvm.nvvm.shfl.sync.down.i32p(i32, i32, i32, i32)297 298declare { float, i1 } @llvm.nvvm.shfl.sync.down.f32p(i32, float, i32, i32)299 300declare { i32, i1 } @llvm.nvvm.shfl.sync.idx.i32p(i32, i32, i32, i32)301 302declare { float, i1 } @llvm.nvvm.shfl.sync.idx.f32p(i32, float, i32, i32)303 304declare i32 @llvm.nvvm.vote.ballot.sync(i32, i1)305 306declare { float, float, float, float, float, float, float, float } @llvm.nvvm.mma.m8n8k4.row.col.f32.f32(<2 x half>, <2 x half>, <2 x half>, <2 x half>, float, float, float, float, float, float, float, float)307 308declare { <2 x half>, <2 x half> } @llvm.nvvm.mma.m16n8k16.row.col.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>)309 310declare { float, float, float, float } @llvm.nvvm.mma.m16n8k16.row.col.f32.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>)311 312declare { <2 x half>, <2 x half> } @llvm.nvvm.mma.m16n8k16.row.col.f16.f32(<2 x half>, <2 x half>, <2 x half>, <2 x half>, <2 x half>, <2 x half>, float, float, float, float)313 314declare { float, float, float, float } @llvm.nvvm.mma.m16n8k16.row.col.f32.f32(<2 x half>, <2 x half>, <2 x half>, <2 x half>, <2 x half>, <2 x half>, float, float, float, float)315 316declare { i32, i32, i32, i32 } @llvm.nvvm.mma.m16n8k16.row.col.s8(i32, i32, i32, i32, i32, i32, i32)317 318declare { i32, i32, i32, i32 } @llvm.nvvm.mma.m16n8k16.row.col.satfinite.s8.u8(i32, i32, i32, i32, i32, i32, i32)319 320declare { i32, i32, i32, i32 } @llvm.nvvm.mma.xor.popc.m16n8k128.row.col.b1(i32, i32, i32, i32, i32, i32, i32)321 322declare { i32, i32, i32, i32 } @llvm.nvvm.mma.m16n8k32.row.col.satfinite.s4(i32, i32, i32, i32, i32, i32, i32)323 324declare { double, double } @llvm.nvvm.mma.m8n8k4.row.col.f64(double, double, double, double)325 326declare { float, float, float, float } @llvm.nvvm.mma.m16n8k4.row.col.tf32(i32, i32, i32, float, float, float, float)327 328declare { <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) nocapture readonly, i32)329 330declare void @llvm.nvvm.wmma.m16n16k16.store.d.row.stride.f16.p3(ptr addrspace(3) nocapture writeonly, <2 x half>, <2 x half>, <2 x half>, <2 x half>, i32)331 332declare { <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>)333 334declare { i32, i32, i32, i32 } @llvm.nvvm.wmma.m16n16k8.load.a.row.stride.tf32.p0(ptr nocapture readonly, i32)335 336declare { 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)337 338declare void @llvm.nvvm.cp.async.ca.shared.global.4(ptr addrspace(3) noalias writeonly, ptr addrspace(1) noalias readonly)339 340declare void @llvm.nvvm.cp.async.ca.shared.global.8(ptr addrspace(3) noalias writeonly, ptr addrspace(1) noalias readonly)341 342declare void @llvm.nvvm.cp.async.ca.shared.global.16(ptr addrspace(3) noalias writeonly, ptr addrspace(1) noalias readonly)343 344declare void @llvm.nvvm.cp.async.cg.shared.global.16(ptr addrspace(3) noalias writeonly, ptr addrspace(1) noalias readonly)345 346declare void @llvm.nvvm.cp.async.commit.group()347 348declare void @llvm.nvvm.cp.async.wait.group(i32 immarg)349 350declare i32 @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x1.b16.p3(ptr addrspace(3) nocapture readonly)351 352declare { i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x2.b16.p3(ptr addrspace(3) nocapture readonly)353 354declare { i32, i32, i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x4.b16.p3(ptr addrspace(3) nocapture readonly)355 356declare i32 @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x1.trans.b16.p3(ptr addrspace(3) nocapture readonly)357 358declare { i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x2.trans.b16.p3(ptr addrspace(3) nocapture readonly)359 360declare { i32, i32, i32, i32 } @llvm.nvvm.ldmatrix.sync.aligned.m8n8.x4.trans.b16.p3(ptr addrspace(3) nocapture readonly)361