279 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 62; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx81 | FileCheck %s3; RUN: %if ptxas-sm_80 && ptxas-isa-8.1 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx81 | %ptxas-verify -arch=sm_80 %}4 5; CHECK-LABEL: cvt_rna_satfinite_tf32_f326define i32 @cvt_rna_satfinite_tf32_f32(float %f1) {7; CHECK-LABEL: cvt_rna_satfinite_tf32_f32(8; CHECK: {9; CHECK-NEXT: .reg .b32 %r<3>;10; CHECK-EMPTY:11; CHECK-NEXT: // %bb.0:12; CHECK-NEXT: ld.param.b32 %r1, [cvt_rna_satfinite_tf32_f32_param_0];13; CHECK-NEXT: cvt.rna.satfinite.tf32.f32 %r2, %r1;14; CHECK-NEXT: st.param.b32 [func_retval0], %r2;15; CHECK-NEXT: ret;16 %val = call i32 @llvm.nvvm.f2tf32.rna.satfinite(float %f1)17 ret i32 %val18}19 20define <2 x bfloat> @cvt_rn_bf16x2_f32_sf(float %f1, float %f2) {21; CHECK-LABEL: cvt_rn_bf16x2_f32_sf(22; CHECK: {23; CHECK-NEXT: .reg .b32 %r<4>;24; CHECK-EMPTY:25; CHECK-NEXT: // %bb.0:26; CHECK-NEXT: ld.param.b32 %r1, [cvt_rn_bf16x2_f32_sf_param_0];27; CHECK-NEXT: ld.param.b32 %r2, [cvt_rn_bf16x2_f32_sf_param_1];28; CHECK-NEXT: cvt.rn.satfinite.bf16x2.f32 %r3, %r1, %r2;29; CHECK-NEXT: st.param.b32 [func_retval0], %r3;30; CHECK-NEXT: ret;31 %val = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rn.satfinite(float %f1, float %f2)32 ret <2 x bfloat> %val33}34 35define <2 x bfloat> @cvt_rn_relu_bf16x2_f32_sf(float %f1, float %f2) {36; CHECK-LABEL: cvt_rn_relu_bf16x2_f32_sf(37; CHECK: {38; CHECK-NEXT: .reg .b32 %r<4>;39; CHECK-EMPTY:40; CHECK-NEXT: // %bb.0:41; CHECK-NEXT: ld.param.b32 %r1, [cvt_rn_relu_bf16x2_f32_sf_param_0];42; CHECK-NEXT: ld.param.b32 %r2, [cvt_rn_relu_bf16x2_f32_sf_param_1];43; CHECK-NEXT: cvt.rn.relu.satfinite.bf16x2.f32 %r3, %r1, %r2;44; CHECK-NEXT: st.param.b32 [func_retval0], %r3;45; CHECK-NEXT: ret;46 %val = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rn.relu.satfinite(float %f1, float %f2)47 ret <2 x bfloat> %val48}49 50define <2 x bfloat> @cvt_rz_bf16x2_f32_sf(float %f1, float %f2) {51; CHECK-LABEL: cvt_rz_bf16x2_f32_sf(52; CHECK: {53; CHECK-NEXT: .reg .b32 %r<4>;54; CHECK-EMPTY:55; CHECK-NEXT: // %bb.0:56; CHECK-NEXT: ld.param.b32 %r1, [cvt_rz_bf16x2_f32_sf_param_0];57; CHECK-NEXT: ld.param.b32 %r2, [cvt_rz_bf16x2_f32_sf_param_1];58; CHECK-NEXT: cvt.rz.satfinite.bf16x2.f32 %r3, %r1, %r2;59; CHECK-NEXT: st.param.b32 [func_retval0], %r3;60; CHECK-NEXT: ret;61 %val = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rz.satfinite(float %f1, float %f2)62 ret <2 x bfloat> %val63}64 65define <2 x bfloat> @cvt_rz_relu_bf16x2_f32_sf(float %f1, float %f2) {66; CHECK-LABEL: cvt_rz_relu_bf16x2_f32_sf(67; CHECK: {68; CHECK-NEXT: .reg .b32 %r<4>;69; CHECK-EMPTY:70; CHECK-NEXT: // %bb.0:71; CHECK-NEXT: ld.param.b32 %r1, [cvt_rz_relu_bf16x2_f32_sf_param_0];72; CHECK-NEXT: ld.param.b32 %r2, [cvt_rz_relu_bf16x2_f32_sf_param_1];73; CHECK-NEXT: cvt.rz.relu.satfinite.bf16x2.f32 %r3, %r1, %r2;74; CHECK-NEXT: st.param.b32 [func_retval0], %r3;75; CHECK-NEXT: ret;76 %val = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rz.relu.satfinite(float %f1, float %f2)77 ret <2 x bfloat> %val78}79 80declare <2 x bfloat> @llvm.nvvm.ff2bf16x2.rn.satfinite(float, float)81declare <2 x bfloat> @llvm.nvvm.ff2bf16x2.rn.relu.satfinite(float, float)82declare <2 x bfloat> @llvm.nvvm.ff2bf16x2.rz.satfinite(float, float)83declare <2 x bfloat> @llvm.nvvm.ff2bf16x2.rz.relu.satfinite(float, float)84 85define <2 x half> @cvt_rn_f16x2_f32_sf(float %f1, float %f2) {86; CHECK-LABEL: cvt_rn_f16x2_f32_sf(87; CHECK: {88; CHECK-NEXT: .reg .b32 %r<4>;89; CHECK-EMPTY:90; CHECK-NEXT: // %bb.0:91; CHECK-NEXT: ld.param.b32 %r1, [cvt_rn_f16x2_f32_sf_param_0];92; CHECK-NEXT: ld.param.b32 %r2, [cvt_rn_f16x2_f32_sf_param_1];93; CHECK-NEXT: cvt.rn.satfinite.f16x2.f32 %r3, %r1, %r2;94; CHECK-NEXT: st.param.b32 [func_retval0], %r3;95; CHECK-NEXT: ret;96 %val = call <2 x half> @llvm.nvvm.ff2f16x2.rn.satfinite(float %f1, float %f2)97 ret <2 x half> %val98}99 100define <2 x half> @cvt_rn_relu_f16x2_f32_sf(float %f1, float %f2) {101; CHECK-LABEL: cvt_rn_relu_f16x2_f32_sf(102; CHECK: {103; CHECK-NEXT: .reg .b32 %r<4>;104; CHECK-EMPTY:105; CHECK-NEXT: // %bb.0:106; CHECK-NEXT: ld.param.b32 %r1, [cvt_rn_relu_f16x2_f32_sf_param_0];107; CHECK-NEXT: ld.param.b32 %r2, [cvt_rn_relu_f16x2_f32_sf_param_1];108; CHECK-NEXT: cvt.rn.relu.satfinite.f16x2.f32 %r3, %r1, %r2;109; CHECK-NEXT: st.param.b32 [func_retval0], %r3;110; CHECK-NEXT: ret;111 %val = call <2 x half> @llvm.nvvm.ff2f16x2.rn.relu.satfinite(float %f1, float %f2)112 ret <2 x half> %val113}114 115define <2 x half> @cvt_rz_f16x2_f32_sf(float %f1, float %f2) {116; CHECK-LABEL: cvt_rz_f16x2_f32_sf(117; CHECK: {118; CHECK-NEXT: .reg .b32 %r<4>;119; CHECK-EMPTY:120; CHECK-NEXT: // %bb.0:121; CHECK-NEXT: ld.param.b32 %r1, [cvt_rz_f16x2_f32_sf_param_0];122; CHECK-NEXT: ld.param.b32 %r2, [cvt_rz_f16x2_f32_sf_param_1];123; CHECK-NEXT: cvt.rz.satfinite.f16x2.f32 %r3, %r1, %r2;124; CHECK-NEXT: st.param.b32 [func_retval0], %r3;125; CHECK-NEXT: ret;126 %val = call <2 x half> @llvm.nvvm.ff2f16x2.rz.satfinite(float %f1, float %f2)127 ret <2 x half> %val128}129 130define <2 x half> @cvt_rz_relu_f16x2_f32_sf(float %f1, float %f2) {131; CHECK-LABEL: cvt_rz_relu_f16x2_f32_sf(132; CHECK: {133; CHECK-NEXT: .reg .b32 %r<4>;134; CHECK-EMPTY:135; CHECK-NEXT: // %bb.0:136; CHECK-NEXT: ld.param.b32 %r1, [cvt_rz_relu_f16x2_f32_sf_param_0];137; CHECK-NEXT: ld.param.b32 %r2, [cvt_rz_relu_f16x2_f32_sf_param_1];138; CHECK-NEXT: cvt.rz.relu.satfinite.f16x2.f32 %r3, %r1, %r2;139; CHECK-NEXT: st.param.b32 [func_retval0], %r3;140; CHECK-NEXT: ret;141 %val = call <2 x half> @llvm.nvvm.ff2f16x2.rz.relu.satfinite(float %f1, float %f2)142 ret <2 x half> %val143}144 145declare <2 x half> @llvm.nvvm.ff2f16x2.rn.satfinite(float, float)146declare <2 x half> @llvm.nvvm.ff2f16x2.rn.relu.satfinite(float, float)147declare <2 x half> @llvm.nvvm.ff2f16x2.rz.satfinite(float, float)148declare <2 x half> @llvm.nvvm.ff2f16x2.rz.relu.satfinite(float, float)149 150define bfloat @cvt_rn_bf16_f32_sf(float %f1) {151; CHECK-LABEL: cvt_rn_bf16_f32_sf(152; CHECK: {153; CHECK-NEXT: .reg .b16 %rs<2>;154; CHECK-NEXT: .reg .b32 %r<2>;155; CHECK-EMPTY:156; CHECK-NEXT: // %bb.0:157; CHECK-NEXT: ld.param.b32 %r1, [cvt_rn_bf16_f32_sf_param_0];158; CHECK-NEXT: cvt.rn.satfinite.bf16.f32 %rs1, %r1;159; CHECK-NEXT: st.param.b16 [func_retval0], %rs1;160; CHECK-NEXT: ret;161 %val = call bfloat @llvm.nvvm.f2bf16.rn.satfinite(float %f1)162 ret bfloat %val163}164 165define bfloat @cvt_rn_relu_bf16_f32_sf(float %f1) {166; CHECK-LABEL: cvt_rn_relu_bf16_f32_sf(167; CHECK: {168; CHECK-NEXT: .reg .b16 %rs<2>;169; CHECK-NEXT: .reg .b32 %r<2>;170; CHECK-EMPTY:171; CHECK-NEXT: // %bb.0:172; CHECK-NEXT: ld.param.b32 %r1, [cvt_rn_relu_bf16_f32_sf_param_0];173; CHECK-NEXT: cvt.rn.relu.satfinite.bf16.f32 %rs1, %r1;174; CHECK-NEXT: st.param.b16 [func_retval0], %rs1;175; CHECK-NEXT: ret;176 %val = call bfloat @llvm.nvvm.f2bf16.rn.relu.satfinite(float %f1)177 ret bfloat %val178}179 180define bfloat @cvt_rz_bf16_f32_sf(float %f1) {181; CHECK-LABEL: cvt_rz_bf16_f32_sf(182; CHECK: {183; CHECK-NEXT: .reg .b16 %rs<2>;184; CHECK-NEXT: .reg .b32 %r<2>;185; CHECK-EMPTY:186; CHECK-NEXT: // %bb.0:187; CHECK-NEXT: ld.param.b32 %r1, [cvt_rz_bf16_f32_sf_param_0];188; CHECK-NEXT: cvt.rz.satfinite.bf16.f32 %rs1, %r1;189; CHECK-NEXT: st.param.b16 [func_retval0], %rs1;190; CHECK-NEXT: ret;191 %val = call bfloat @llvm.nvvm.f2bf16.rz.satfinite(float %f1)192 ret bfloat %val193}194 195define bfloat @cvt_rz_relu_bf16_f32_sf(float %f1) {196; CHECK-LABEL: cvt_rz_relu_bf16_f32_sf(197; CHECK: {198; CHECK-NEXT: .reg .b16 %rs<2>;199; CHECK-NEXT: .reg .b32 %r<2>;200; CHECK-EMPTY:201; CHECK-NEXT: // %bb.0:202; CHECK-NEXT: ld.param.b32 %r1, [cvt_rz_relu_bf16_f32_sf_param_0];203; CHECK-NEXT: cvt.rz.relu.satfinite.bf16.f32 %rs1, %r1;204; CHECK-NEXT: st.param.b16 [func_retval0], %rs1;205; CHECK-NEXT: ret;206 %val = call bfloat @llvm.nvvm.f2bf16.rz.relu.satfinite(float %f1)207 ret bfloat %val208}209 210declare bfloat @llvm.nvvm.f2bf16.rn.satfinite(float)211declare bfloat @llvm.nvvm.f2bf16.rn.relu.satfinite(float)212declare bfloat @llvm.nvvm.f2bf16.rz.satfinite(float)213declare bfloat @llvm.nvvm.f2bf16.rz.relu.satfinite(float)214 215define half @cvt_rn_f16_f32_sf(float %f1) {216; CHECK-LABEL: cvt_rn_f16_f32_sf(217; CHECK: {218; CHECK-NEXT: .reg .b16 %rs<2>;219; CHECK-NEXT: .reg .b32 %r<2>;220; CHECK-EMPTY:221; CHECK-NEXT: // %bb.0:222; CHECK-NEXT: ld.param.b32 %r1, [cvt_rn_f16_f32_sf_param_0];223; CHECK-NEXT: cvt.rn.satfinite.f16.f32 %rs1, %r1;224; CHECK-NEXT: st.param.b16 [func_retval0], %rs1;225; CHECK-NEXT: ret;226 %val = call half @llvm.nvvm.f2f16.rn.satfinite(float %f1)227 ret half %val228}229 230define half @cvt_rn_relu_f16_f32_sf(float %f1) {231; CHECK-LABEL: cvt_rn_relu_f16_f32_sf(232; CHECK: {233; CHECK-NEXT: .reg .b16 %rs<2>;234; CHECK-NEXT: .reg .b32 %r<2>;235; CHECK-EMPTY:236; CHECK-NEXT: // %bb.0:237; CHECK-NEXT: ld.param.b32 %r1, [cvt_rn_relu_f16_f32_sf_param_0];238; CHECK-NEXT: cvt.rn.relu.satfinite.f16.f32 %rs1, %r1;239; CHECK-NEXT: st.param.b16 [func_retval0], %rs1;240; CHECK-NEXT: ret;241 %val = call half @llvm.nvvm.f2f16.rn.relu.satfinite(float %f1)242 ret half %val243}244 245define half @cvt_rz_f16_f32_sf(float %f1) {246; CHECK-LABEL: cvt_rz_f16_f32_sf(247; CHECK: {248; CHECK-NEXT: .reg .b16 %rs<2>;249; CHECK-NEXT: .reg .b32 %r<2>;250; CHECK-EMPTY:251; CHECK-NEXT: // %bb.0:252; CHECK-NEXT: ld.param.b32 %r1, [cvt_rz_f16_f32_sf_param_0];253; CHECK-NEXT: cvt.rz.satfinite.f16.f32 %rs1, %r1;254; CHECK-NEXT: st.param.b16 [func_retval0], %rs1;255; CHECK-NEXT: ret;256 %val = call half @llvm.nvvm.f2f16.rz.satfinite(float %f1)257 ret half %val258}259 260define half @cvt_rz_relu_f16_f32_sf(float %f1) {261; CHECK-LABEL: cvt_rz_relu_f16_f32_sf(262; CHECK: {263; CHECK-NEXT: .reg .b16 %rs<2>;264; CHECK-NEXT: .reg .b32 %r<2>;265; CHECK-EMPTY:266; CHECK-NEXT: // %bb.0:267; CHECK-NEXT: ld.param.b32 %r1, [cvt_rz_relu_f16_f32_sf_param_0];268; CHECK-NEXT: cvt.rz.relu.satfinite.f16.f32 %rs1, %r1;269; CHECK-NEXT: st.param.b16 [func_retval0], %rs1;270; CHECK-NEXT: ret;271 %val = call half @llvm.nvvm.f2f16.rz.relu.satfinite(float %f1)272 ret half %val273}274 275declare half @llvm.nvvm.f2f16.rn.satfinite(float)276declare half @llvm.nvvm.f2f16.rn.relu.satfinite(float)277declare half @llvm.nvvm.f2f16.rz.satfinite(float)278declare half @llvm.nvvm.f2f16.rz.relu.satfinite(float)279