375 lines · plain
1; Test to make sure NVVM intrinsics are automatically upgraded.2; RUN: llvm-as < %s | llvm-dis | FileCheck %s3; RUN: verify-uselistorder %s4 5declare i32 @llvm.nvvm.brev32(i32)6declare i64 @llvm.nvvm.brev64(i64)7declare i32 @llvm.nvvm.clz.i(i32)8declare i32 @llvm.nvvm.clz.ll(i64)9declare i32 @llvm.nvvm.popc.i(i32)10declare i32 @llvm.nvvm.popc.ll(i64)11declare float @llvm.nvvm.h2f(i16)12 13declare i32 @llvm.nvvm.abs.i(i32)14declare i64 @llvm.nvvm.abs.ll(i64)15 16declare float @llvm.nvvm.fabs.f(float)17declare float @llvm.nvvm.fabs.ftz.f(float)18declare double @llvm.nvvm.fabs.d(double)19 20declare float @llvm.nvvm.tanh.approx.f32(float)21 22declare i16 @llvm.nvvm.max.s(i16, i16)23declare i32 @llvm.nvvm.max.i(i32, i32)24declare i64 @llvm.nvvm.max.ll(i64, i64)25declare i16 @llvm.nvvm.max.us(i16, i16)26declare i32 @llvm.nvvm.max.ui(i32, i32)27declare i64 @llvm.nvvm.max.ull(i64, i64)28declare i16 @llvm.nvvm.min.s(i16, i16)29declare i32 @llvm.nvvm.min.i(i32, i32)30declare i64 @llvm.nvvm.min.ll(i64, i64)31declare i16 @llvm.nvvm.min.us(i16, i16)32declare i32 @llvm.nvvm.min.ui(i32, i32)33declare i64 @llvm.nvvm.min.ull(i64, i64)34 35declare i32 @llvm.nvvm.bitcast.f2i(float)36declare float @llvm.nvvm.bitcast.i2f(i32)37declare i64 @llvm.nvvm.bitcast.d2ll(double)38declare double @llvm.nvvm.bitcast.ll2d(i64)39 40declare i32 @llvm.nvvm.rotate.b32(i32, i32)41declare i64 @llvm.nvvm.rotate.right.b64(i64, i32)42declare i64 @llvm.nvvm.rotate.b64(i64, i32)43declare i64 @llvm.nvvm.swap.lo.hi.b64(i64)44 45declare ptr addrspace(1) @llvm.nvvm.ptr.gen.to.global.p1.p0(ptr)46declare ptr addrspace(3) @llvm.nvvm.ptr.gen.to.shared.p3.p0(ptr)47declare ptr addrspace(4) @llvm.nvvm.ptr.gen.to.constant.p4.p0(ptr)48declare ptr addrspace(5) @llvm.nvvm.ptr.gen.to.local.p5.p0(ptr)49declare ptr addrspace(101) @llvm.nvvm.ptr.gen.to.param.p101.p0(ptr)50declare ptr @llvm.nvvm.ptr.global.to.gen.p0.p1(ptr addrspace(1))51declare ptr @llvm.nvvm.ptr.shared.to.gen.p0.p3(ptr addrspace(3))52declare ptr @llvm.nvvm.ptr.constant.to.gen.p0.p4(ptr addrspace(4))53declare ptr @llvm.nvvm.ptr.local.to.gen.p0.p5(ptr addrspace(5))54declare ptr @llvm.nvvm.ptr.param.to.gen.p0.p101(ptr addrspace(101))55 56declare i32 @llvm.nvvm.ldg.global.i.i32.p1(ptr addrspace(1), i32)57declare ptr @llvm.nvvm.ldg.global.p.p1(ptr addrspace(1), i32)58declare float @llvm.nvvm.ldg.global.f.f32.p1(ptr addrspace(1), i32)59declare i32 @llvm.nvvm.ldg.global.i.i32.p0(ptr, i32)60declare ptr @llvm.nvvm.ldg.global.p.p0(ptr, i32)61declare float @llvm.nvvm.ldg.global.f.f32.p0(ptr, i32)62 63declare i32 @llvm.nvvm.atomic.load.inc.32.p0(ptr, i32)64declare i32 @llvm.nvvm.atomic.load.dec.32.p0(ptr, i32)65declare i32 @llvm.nvvm.atomic.load.add.f32.p0(ptr, float)66declare i32 @llvm.nvvm.atomic.load.add.f64.p0(ptr, double)67 68declare ptr addrspace(3) @llvm.nvvm.mapa.shared.cluster(ptr addrspace(3), i32)69 70declare void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster(ptr addrspace(3), ptr addrspace(3), ptr addrspace(1), i32, i16, i64, i1, i1)71declare void @llvm.nvvm.cp.async.bulk.shared.cta.to.cluster(ptr addrspace(3), ptr addrspace(3), ptr addrspace(3), i32)72 73declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.1d(ptr addrspace(3) %d, ptr addrspace(3) %bar, ptr %tm, i32 %d0, i16 %mc, i64 %ch, i1 %f1, i1 %f2);74declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.2d(ptr addrspace(3) %d, ptr addrspace(3) %bar, ptr %tm, i32 %d0, i32 %d1, i16 %mc, i64 %ch, i1 %f1, i1 %f2);75declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.3d(ptr addrspace(3) %d, ptr addrspace(3) %bar, ptr %tm, i32 %d0, i32 %d1, i32 %d2, i16 %mc, i64 %ch, i1 %f1, i1 %f2);76declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.4d(ptr addrspace(3) %d, ptr addrspace(3) %bar, ptr %tm, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i16 %mc, i64 %ch, i1 %f1, i1 %f2);77declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.5d(ptr addrspace(3) %d, ptr addrspace(3) %bar, ptr %tm, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, i16 %mc, i64 %ch, i1 %f1, i1 %f2);78 79declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.3d(ptr addrspace(3) %d, ptr addrspace(3) %bar, ptr %tm, i32 %d0, i32 %d1, i32 %d2, i16 %im2col0, i16 %mc, i64 %ch, i1 %f1, i1 %f2);80declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.4d(ptr addrspace(3) %d, ptr addrspace(3) %bar, ptr %tm, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i16 %im2col0, i16 %im2col1, i16 %mc, i64 %ch, i1 %f1, i1 %f2);81declare void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.5d(ptr addrspace(3) %d, ptr addrspace(3) %bar, ptr %tm, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, i16 %im2col0, i16 %im2col1, i16 %im2col2, i16 %mc, i64 %ch, i1 %f1, i1 %f2);82 83declare void @llvm.nvvm.barrier0()84declare void @llvm.nvvm.barrier.n(i32)85declare void @llvm.nvvm.bar.sync(i32)86declare void @llvm.nvvm.barrier(i32, i32)87declare void @llvm.nvvm.barrier.sync(i32)88declare void @llvm.nvvm.barrier.sync.cnt(i32, i32)89 90declare float @llvm.nvvm.ex2.approx.f(float)91declare double @llvm.nvvm.ex2.approx.d(double)92declare <2 x half> @llvm.nvvm.ex2.approx.f16x2(<2 x half>)93declare float @llvm.nvvm.ex2.approx.ftz.f(float)94 95; CHECK-LABEL: @simple_upgrade96define void @simple_upgrade(i32 %a, i64 %b, i16 %c) {97; CHECK: call i32 @llvm.bitreverse.i32(i32 %a)98 %r1 = call i32 @llvm.nvvm.brev32(i32 %a)99 100; CHECK: call i64 @llvm.bitreverse.i64(i64 %b)101 %r2 = call i64 @llvm.nvvm.brev64(i64 %b)102 103; CHECK: call i32 @llvm.ctlz.i32(i32 %a, i1 false)104 %r3 = call i32 @llvm.nvvm.clz.i(i32 %a)105 106; CHECK: [[clz:%[a-zA-Z0-9.]+]] = call i64 @llvm.ctlz.i64(i64 %b, i1 false)107; CHECK: trunc i64 [[clz]] to i32108 %r4 = call i32 @llvm.nvvm.clz.ll(i64 %b)109 110; CHECK: call i32 @llvm.ctpop.i32(i32 %a)111 %r5 = call i32 @llvm.nvvm.popc.i(i32 %a)112 113; CHECK: [[popc:%[a-zA-Z0-9.]+]] = call i64 @llvm.ctpop.i64(i64 %b)114; CHECK: trunc i64 [[popc]] to i32115 %r6 = call i32 @llvm.nvvm.popc.ll(i64 %b)116 117; CHECK: call float @llvm.convert.from.fp16.f32(i16 %c)118 %r7 = call float @llvm.nvvm.h2f(i16 %c)119 ret void120}121 122; CHECK-LABEL: @abs123define void @abs(i32 %a, i64 %b) {124; CHECK-DAG: [[negi:%[a-zA-Z0-9.]+]] = sub i32 0, %a125; CHECK-DAG: [[cmpi:%[a-zA-Z0-9.]+]] = icmp sge i32 %a, 0126; CHECK: select i1 [[cmpi]], i32 %a, i32 [[negi]]127 %r1 = call i32 @llvm.nvvm.abs.i(i32 %a)128 129; CHECK-DAG: [[negll:%[a-zA-Z0-9.]+]] = sub i64 0, %b130; CHECK-DAG: [[cmpll:%[a-zA-Z0-9.]+]] = icmp sge i64 %b, 0131; CHECK: select i1 [[cmpll]], i64 %b, i64 [[negll]]132 %r2 = call i64 @llvm.nvvm.abs.ll(i64 %b)133 134 ret void135}136 137; CHECK-LABEL: @fabs138define void @fabs(float %a, double %b) {139; CHECK: call float @llvm.nvvm.fabs.f32(float %a)140; CHECK: call float @llvm.nvvm.fabs.ftz.f32(float %a)141; CHECK: call double @llvm.nvvm.fabs.f64(double %b)142 %r1 = call float @llvm.nvvm.fabs.f(float %a)143 %r2 = call float @llvm.nvvm.fabs.ftz.f(float %a)144 %r3 = call double @llvm.nvvm.fabs.d(double %b)145 ret void146}147 148; CHECK-LABEL: @tanh149define void @tanh(float %a) {150; CHECK: call afn float @llvm.tanh.f32(float %a)151 %r1 = call float @llvm.nvvm.tanh.approx.f32(float %a)152 ret void153}154 155; CHECK-LABEL: @min_max156define void @min_max(i16 %a1, i16 %a2, i32 %b1, i32 %b2, i64 %c1, i64 %c2) {157; CHECK: [[maxs:%[a-zA-Z0-9.]+]] = icmp sge i16 %a1, %a2158; CHECK: select i1 [[maxs]], i16 %a1, i16 %a2159 %r1 = call i16 @llvm.nvvm.max.s(i16 %a1, i16 %a2)160 161; CHECK: [[maxi:%[a-zA-Z0-9.]+]] = icmp sge i32 %b1, %b2162; CHECK: select i1 [[maxi]], i32 %b1, i32 %b2163 %r2 = call i32 @llvm.nvvm.max.i(i32 %b1, i32 %b2)164 165; CHECK: [[maxll:%[a-zA-Z0-9.]+]] = icmp sge i64 %c1, %c2166; CHECK: select i1 [[maxll]], i64 %c1, i64 %c2167 %r3 = call i64 @llvm.nvvm.max.ll(i64 %c1, i64 %c2)168 169; CHECK: [[maxus:%[a-zA-Z0-9.]+]] = icmp uge i16 %a1, %a2170; CHECK: select i1 [[maxus]], i16 %a1, i16 %a2171 %r4 = call i16 @llvm.nvvm.max.us(i16 %a1, i16 %a2)172 173; CHECK: [[maxui:%[a-zA-Z0-9.]+]] = icmp uge i32 %b1, %b2174; CHECK: select i1 [[maxui]], i32 %b1, i32 %b2175 %r5 = call i32 @llvm.nvvm.max.ui(i32 %b1, i32 %b2)176 177; CHECK: [[maxull:%[a-zA-Z0-9.]+]] = icmp uge i64 %c1, %c2178; CHECK: select i1 [[maxull]], i64 %c1, i64 %c2179 %r6 = call i64 @llvm.nvvm.max.ull(i64 %c1, i64 %c2)180 181; CHECK: [[mins:%[a-zA-Z0-9.]+]] = icmp sle i16 %a1, %a2182; CHECK: select i1 [[mins]], i16 %a1, i16 %a2183 %r7 = call i16 @llvm.nvvm.min.s(i16 %a1, i16 %a2)184 185; CHECK: [[mini:%[a-zA-Z0-9.]+]] = icmp sle i32 %b1, %b2186; CHECK: select i1 [[mini]], i32 %b1, i32 %b2187 %r8 = call i32 @llvm.nvvm.min.i(i32 %b1, i32 %b2)188 189; CHECK: [[minll:%[a-zA-Z0-9.]+]] = icmp sle i64 %c1, %c2190; CHECK: select i1 [[minll]], i64 %c1, i64 %c2191 %r9 = call i64 @llvm.nvvm.min.ll(i64 %c1, i64 %c2)192 193; CHECK: [[minus:%[a-zA-Z0-9.]+]] = icmp ule i16 %a1, %a2194; CHECK: select i1 [[minus]], i16 %a1, i16 %a2195 %r10 = call i16 @llvm.nvvm.min.us(i16 %a1, i16 %a2)196 197; CHECK: [[minui:%[a-zA-Z0-9.]+]] = icmp ule i32 %b1, %b2198; CHECK: select i1 [[minui]], i32 %b1, i32 %b2199 %r11 = call i32 @llvm.nvvm.min.ui(i32 %b1, i32 %b2)200 201; CHECK: [[minull:%[a-zA-Z0-9.]+]] = icmp ule i64 %c1, %c2202; CHECK: select i1 [[minull]], i64 %c1, i64 %c2203 %r12 = call i64 @llvm.nvvm.min.ull(i64 %c1, i64 %c2)204 205 ret void206}207 208; CHECK-LABEL: @bitcast209define void @bitcast(i32 %a, i64 %b, float %c, double %d) {210; CHECK: bitcast float %c to i32211; CHECK: bitcast i32 %a to float212; CHECK: bitcast double %d to i64213; CHECK: bitcast i64 %b to double214;215 %r1 = call i32 @llvm.nvvm.bitcast.f2i(float %c)216 %r2 = call float @llvm.nvvm.bitcast.i2f(i32 %a)217 %r3 = call i64 @llvm.nvvm.bitcast.d2ll(double %d)218 %r4 = call double @llvm.nvvm.bitcast.ll2d(i64 %b)219 220 ret void221}222 223; CHECK-LABEL: @rotate224define void @rotate(i32 %a, i64 %b) {225; CHECK: call i32 @llvm.fshl.i32(i32 %a, i32 %a, i32 6)226; CHECK: call i64 @llvm.fshr.i64(i64 %b, i64 %b, i64 7)227; CHECK: call i64 @llvm.fshl.i64(i64 %b, i64 %b, i64 8)228; CHECK: call i64 @llvm.fshl.i64(i64 %b, i64 %b, i64 32)229;230 %r1 = call i32 @llvm.nvvm.rotate.b32(i32 %a, i32 6)231 %r2 = call i64 @llvm.nvvm.rotate.right.b64(i64 %b, i32 7)232 %r3 = call i64 @llvm.nvvm.rotate.b64(i64 %b, i32 8)233 %r4 = call i64 @llvm.nvvm.swap.lo.hi.b64(i64 %b)234 ret void235}236 237; CHECK-LABEL: @addrspacecast238define void @addrspacecast(ptr %p0) {239; CHECK: %1 = addrspacecast ptr %p0 to ptr addrspace(1)240; CHECK: %2 = addrspacecast ptr addrspace(1) %1 to ptr241; CHECK: %3 = addrspacecast ptr %2 to ptr addrspace(3)242; CHECK: %4 = addrspacecast ptr addrspace(3) %3 to ptr243; CHECK: %5 = addrspacecast ptr %4 to ptr addrspace(4)244; CHECK: %6 = addrspacecast ptr addrspace(4) %5 to ptr245; CHECK: %7 = addrspacecast ptr %6 to ptr addrspace(5)246; CHECK: %8 = addrspacecast ptr addrspace(5) %7 to ptr247; CHECK: %9 = addrspacecast ptr %8 to ptr addrspace(101)248; CHECK: %10 = addrspacecast ptr addrspace(101) %9 to ptr249;250 %p1 = call ptr addrspace(1) @llvm.nvvm.ptr.gen.to.global.p1.p0(ptr %p0)251 %p2 = call ptr @llvm.nvvm.ptr.global.to.gen.p0.p1(ptr addrspace(1) %p1)252 253 %p3 = call ptr addrspace(3) @llvm.nvvm.ptr.gen.to.shared.p3.p0(ptr %p2)254 %p4 = call ptr @llvm.nvvm.ptr.shared.to.gen.p0.p3(ptr addrspace(3) %p3)255 256 %p5 = call ptr addrspace(4) @llvm.nvvm.ptr.gen.to.constant.p4.p0(ptr %p4)257 %p6 = call ptr @llvm.nvvm.ptr.constant.to.gen.p0.p4(ptr addrspace(4) %p5)258 259 %p7 = call ptr addrspace(5) @llvm.nvvm.ptr.gen.to.local.p5.p0(ptr %p6)260 %p8 = call ptr @llvm.nvvm.ptr.local.to.gen.p0.p5(ptr addrspace(5) %p7)261 262 %p9 = call ptr addrspace(101) @llvm.nvvm.ptr.gen.to.param.p101.p0(ptr %p8)263 %p10 = call ptr @llvm.nvvm.ptr.param.to.gen.p0.p101(ptr addrspace(101) %p9)264 265 ret void266}267 268; CHECK-LABEL: @ldg269define void @ldg(ptr %p0, ptr addrspace(1) %p1) {270; CHECK: %1 = load i32, ptr addrspace(1) %p1, align 4, !invariant.load !0271; CHECK: %2 = load ptr, ptr addrspace(1) %p1, align 8, !invariant.load !0272; CHECK: %3 = load float, ptr addrspace(1) %p1, align 16, !invariant.load !0273 274; CHECK: %4 = addrspacecast ptr %p0 to ptr addrspace(1)275; CHECK: %5 = load i32, ptr addrspace(1) %4, align 4, !invariant.load !0276; CHECK: %6 = addrspacecast ptr %p0 to ptr addrspace(1)277; CHECK: %7 = load ptr, ptr addrspace(1) %6, align 8, !invariant.load !0278; CHECK: %8 = addrspacecast ptr %p0 to ptr addrspace(1)279; CHECK: %9 = load float, ptr addrspace(1) %8, align 16, !invariant.load !0280;281 %v1 = call i32 @llvm.nvvm.ldg.global.i.i32.p1(ptr addrspace(1) %p1, i32 4)282 %v2 = call ptr @llvm.nvvm.ldg.global.p.p1(ptr addrspace(1) %p1, i32 8 )283 %v3 = call float @llvm.nvvm.ldg.global.f.f32.p1(ptr addrspace(1) %p1, i32 16)284 285 %v4 = call i32 @llvm.nvvm.ldg.global.i.i32.p0(ptr %p0, i32 4)286 %v5 = call ptr @llvm.nvvm.ldg.global.p.p0(ptr %p0, i32 8)287 %v6 = call float @llvm.nvvm.ldg.global.f.f32.p0(ptr %p0, i32 16)288 289 ret void290}291 292; CHECK-LABEL: @atomics293define i32 @atomics(ptr %p0, i32 %a, float %b, double %c) {294; CHECK: %1 = atomicrmw uinc_wrap ptr %p0, i32 %a seq_cst295; CHECK: %2 = atomicrmw udec_wrap ptr %p0, i32 %a seq_cst296; CHECK: %3 = atomicrmw fadd ptr %p0, float %b seq_cst297; CHECK: %4 = atomicrmw fadd ptr %p0, double %c seq_cst298 299 %r1 = call i32 @llvm.nvvm.atomic.load.inc.32.p0(ptr %p0, i32 %a)300 %r2 = call i32 @llvm.nvvm.atomic.load.dec.32.p0(ptr %p0, i32 %a)301 %r3 = call float @llvm.nvvm.atomic.load.add.f32.p0(ptr %p0, float %b)302 %r4 = call double @llvm.nvvm.atomic.load.add.f64.p0(ptr %p0, double %c)303 ret i32 %r2304}305 306; CHECK-LABEL: @nvvm_shared_cluster_intrinsics307define void @nvvm_shared_cluster_intrinsics(ptr addrspace(3) %p0, i32 %offset) {308; CHECK: %r = call ptr addrspace(7) @llvm.nvvm.mapa.shared.cluster(ptr addrspace(3) %p0, i32 %offset)309 %r = call ptr addrspace(3) @llvm.nvvm.mapa.shared.cluster(ptr addrspace(3) %p0, i32 %offset)310 ret void311}312 313; CHECK-LABEL: @nvvm_cp_async_bulk_intrinsics314define void @nvvm_cp_async_bulk_intrinsics(ptr addrspace(3) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, ptr addrspace(3) %src_shared, i32 %size) {315; CHECK: call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster(ptr addrspace(7) %1, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i16 0, i64 0, i1 false, i1 false)316; CHECK: call void @llvm.nvvm.cp.async.bulk.shared.cta.to.cluster(ptr addrspace(7) %2, ptr addrspace(3) %bar, ptr addrspace(3) %src_shared, i32 %size)317 call void @llvm.nvvm.cp.async.bulk.global.to.shared.cluster(ptr addrspace(3) %dst, ptr addrspace(3) %bar, ptr addrspace(1) %src, i32 %size, i16 0, i64 0, i1 false, i1 false)318 call void @llvm.nvvm.cp.async.bulk.shared.cta.to.cluster(ptr addrspace(3) %dst, ptr addrspace(3) %bar, ptr addrspace(3) %src_shared, i32 %size)319 ret void320}321 322; CHECK-LABEL: @nvvm_cp_async_bulk_tensor_g2s_im2col323define void @nvvm_cp_async_bulk_tensor_g2s_im2col(ptr addrspace(3) %d, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, i16 %im2col0, i16 %im2col1, i16 %im2col2, i16 %mc, i64 %ch) {324; CHECK: call void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.3d(ptr addrspace(7) %1, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i16 %im2col0, i16 0, i64 0, i1 false, i1 false, i32 0)325; CHECK: call void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.4d(ptr addrspace(7) %2, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i16 %im2col0, i16 %im2col1, i16 0, i64 0, i1 false, i1 false, i32 0)326; CHECK: call void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.5d(ptr addrspace(7) %3, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, i16 %im2col0, i16 %im2col1, i16 %im2col2, i16 0, i64 0, i1 false, i1 false, i32 0)327 call void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.3d(ptr addrspace(3) %d, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i16 %im2col0, i16 0, i64 0, i1 0, i1 0)328 call void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.4d(ptr addrspace(3) %d, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i16 %im2col0, i16 %im2col1, i16 0, i64 0, i1 0, i1 0)329 call void @llvm.nvvm.cp.async.bulk.tensor.g2s.im2col.5d(ptr addrspace(3) %d, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, i16 %im2col0, i16 %im2col1, i16 %im2col2, i16 0, i64 0, i1 0, i1 0)330 ret void331}332 333; CHECK-LABEL: @nvvm_cp_async_bulk_tensor_g2s_tile334define void @nvvm_cp_async_bulk_tensor_g2s_tile(ptr addrspace(3) %d, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, i16 %mc, i64 %ch) {335; CHECK: call void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.1d(ptr addrspace(7) %1, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i16 0, i64 0, i1 false, i1 false, i32 0)336; CHECK: call void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.2d(ptr addrspace(7) %2, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i32 %d1, i16 0, i64 0, i1 false, i1 false, i32 0)337; CHECK: call void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.3d(ptr addrspace(7) %3, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i16 0, i64 0, i1 false, i1 false, i32 0)338; CHECK: call void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.4d(ptr addrspace(7) %4, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i16 0, i64 0, i1 false, i1 false, i32 0)339; CHECK: call void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.5d(ptr addrspace(7) %5, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, i16 0, i64 0, i1 false, i1 false, i32 0)340 call void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.1d(ptr addrspace(3) %d, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i16 0, i64 0, i1 0, i1 0)341 call void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.2d(ptr addrspace(3) %d, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i32 %d1, i16 0, i64 0, i1 0, i1 0)342 call void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.3d(ptr addrspace(3) %d, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i16 0, i64 0, i1 0, i1 0)343 call void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.4d(ptr addrspace(3) %d, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i16 0, i64 0, i1 0, i1 0)344 call void @llvm.nvvm.cp.async.bulk.tensor.g2s.tile.5d(ptr addrspace(3) %d, ptr addrspace(3) %bar, ptr %tmap, i32 %d0, i32 %d1, i32 %d2, i32 %d3, i32 %d4, i16 0, i64 0, i1 0, i1 0)345 ret void346}347 348define void @cta_barriers(i32 %x, i32 %y) {349; CHECK: call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 0)350; CHECK: call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 %x)351; CHECK: call void @llvm.nvvm.barrier.cta.sync.aligned.all(i32 %x)352; CHECK: call void @llvm.nvvm.barrier.cta.sync.aligned.count(i32 %x, i32 %y)353; CHECK: call void @llvm.nvvm.barrier.cta.sync.all(i32 %x)354; CHECK: call void @llvm.nvvm.barrier.cta.sync.count(i32 %x, i32 %y)355 call void @llvm.nvvm.barrier0()356 call void @llvm.nvvm.barrier.n(i32 %x)357 call void @llvm.nvvm.bar.sync(i32 %x)358 call void @llvm.nvvm.barrier(i32 %x, i32 %y)359 call void @llvm.nvvm.barrier.sync(i32 %x)360 call void @llvm.nvvm.barrier.sync.cnt(i32 %x, i32 %y)361 ret void362}363 364define void @nvvm_ex2_approx(float %a, double %b, half %c, <2 x half> %d) {365; CHECK: call float @llvm.nvvm.ex2.approx.f32(float %a)366; CHECK: call double @llvm.nvvm.ex2.approx.f64(double %b)367; CHECK: call <2 x half> @llvm.nvvm.ex2.approx.v2f16(<2 x half> %d)368; CHECK: call float @llvm.nvvm.ex2.approx.ftz.f32(float %a)369 %r1 = call float @llvm.nvvm.ex2.approx.f(float %a)370 %r2 = call double @llvm.nvvm.ex2.approx.d(double %b)371 %r3 = call <2 x half> @llvm.nvvm.ex2.approx.f16x2(<2 x half> %d)372 %r4 = call float @llvm.nvvm.ex2.approx.ftz.f(float %a)373 ret void374}375