66 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mcpu=sm_90 -mattr=+ptx78 | FileCheck --check-prefixes=CHECK-FP16 %s3; RUN: %if ptxas-sm_90 && ptxas-isa-7.8 %{ llc < %s -mcpu=sm_90 -mattr=+ptx78 | %ptxas-verify -arch=sm_90 %}4target triple = "nvptx64-nvidia-cuda"5 6declare half @llvm.nvvm.ex2.approx.f16(half)7declare <2 x half> @llvm.nvvm.ex2.approx.v2f16(<2 x half>)8declare bfloat @llvm.nvvm.ex2.approx.ftz.bf16(bfloat)9declare <2 x bfloat> @llvm.nvvm.ex2.approx.ftz.v2bf16(<2 x bfloat>)10 11define half @ex2_half(half %0) {12; CHECK-FP16-LABEL: ex2_half(13; CHECK-FP16: {14; CHECK-FP16-NEXT: .reg .b16 %rs<3>;15; CHECK-FP16-EMPTY:16; CHECK-FP16-NEXT: // %bb.0:17; CHECK-FP16-NEXT: ld.param.b16 %rs1, [ex2_half_param_0];18; CHECK-FP16-NEXT: ex2.approx.f16 %rs2, %rs1;19; CHECK-FP16-NEXT: st.param.b16 [func_retval0], %rs2;20; CHECK-FP16-NEXT: ret;21 %res = call half @llvm.nvvm.ex2.approx.f16(half %0)22 ret half %res23}24 25define <2 x half> @ex2_2xhalf(<2 x half> %0) {26; CHECK-FP16-LABEL: ex2_2xhalf(27; CHECK-FP16: {28; CHECK-FP16-NEXT: .reg .b32 %r<3>;29; CHECK-FP16-EMPTY:30; CHECK-FP16-NEXT: // %bb.0:31; CHECK-FP16-NEXT: ld.param.b32 %r1, [ex2_2xhalf_param_0];32; CHECK-FP16-NEXT: ex2.approx.f16x2 %r2, %r1;33; CHECK-FP16-NEXT: st.param.b32 [func_retval0], %r2;34; CHECK-FP16-NEXT: ret;35 %res = call <2 x half> @llvm.nvvm.ex2.approx.v2f16(<2 x half> %0)36 ret <2 x half> %res37}38 39define bfloat @ex2_bfloat(bfloat %0) {40; CHECK-FP16-LABEL: ex2_bfloat(41; CHECK-FP16: {42; CHECK-FP16-NEXT: .reg .b16 %rs<3>;43; CHECK-FP16-EMPTY:44; CHECK-FP16-NEXT: // %bb.0:45; CHECK-FP16-NEXT: ld.param.b16 %rs1, [ex2_bfloat_param_0];46; CHECK-FP16-NEXT: ex2.approx.ftz.bf16 %rs2, %rs1;47; CHECK-FP16-NEXT: st.param.b16 [func_retval0], %rs2;48; CHECK-FP16-NEXT: ret;49 %res = call bfloat @llvm.nvvm.ex2.approx.ftz.bf16(bfloat %0)50 ret bfloat %res51}52 53define <2 x bfloat> @ex2_2xbfloat(<2 x bfloat> %0) {54; CHECK-FP16-LABEL: ex2_2xbfloat(55; CHECK-FP16: {56; CHECK-FP16-NEXT: .reg .b32 %r<3>;57; CHECK-FP16-EMPTY:58; CHECK-FP16-NEXT: // %bb.0:59; CHECK-FP16-NEXT: ld.param.b32 %r1, [ex2_2xbfloat_param_0];60; CHECK-FP16-NEXT: ex2.approx.ftz.bf16x2 %r2, %r1;61; CHECK-FP16-NEXT: st.param.b32 [func_retval0], %r2;62; CHECK-FP16-NEXT: ret;63 %res = call <2 x bfloat> @llvm.nvvm.ex2.approx.ftz.v2bf16(<2 x bfloat> %0)64 ret <2 x bfloat> %res65}66