318 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx70 | FileCheck %s3; RUN: %if ptxas-sm_80 && ptxas-isa-7.0 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx70 | %ptxas-verify -arch=sm_80 %}4 5 6define <2 x bfloat> @cvt_rn_bf16x2_f32(float %f1, float %f2) {7; CHECK-LABEL: cvt_rn_bf16x2_f32(8; CHECK: {9; CHECK-NEXT: .reg .b32 %r<4>;10; CHECK-EMPTY:11; CHECK-NEXT: // %bb.0:12; CHECK-NEXT: ld.param.b32 %r1, [cvt_rn_bf16x2_f32_param_0];13; CHECK-NEXT: ld.param.b32 %r2, [cvt_rn_bf16x2_f32_param_1];14; CHECK-NEXT: cvt.rn.bf16x2.f32 %r3, %r1, %r2;15; CHECK-NEXT: st.param.b32 [func_retval0], %r3;16; CHECK-NEXT: ret;17 %val = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rn(float %f1, float %f2)18 ret <2 x bfloat> %val19}20 21define <2 x bfloat> @cvt_rn_relu_bf16x2_f32(float %f1, float %f2) {22; CHECK-LABEL: cvt_rn_relu_bf16x2_f32(23; CHECK: {24; CHECK-NEXT: .reg .b32 %r<4>;25; CHECK-EMPTY:26; CHECK-NEXT: // %bb.0:27; CHECK-NEXT: ld.param.b32 %r1, [cvt_rn_relu_bf16x2_f32_param_0];28; CHECK-NEXT: ld.param.b32 %r2, [cvt_rn_relu_bf16x2_f32_param_1];29; CHECK-NEXT: cvt.rn.relu.bf16x2.f32 %r3, %r1, %r2;30; CHECK-NEXT: st.param.b32 [func_retval0], %r3;31; CHECK-NEXT: ret;32 %val = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rn.relu(float %f1, float %f2)33 ret <2 x bfloat> %val34}35 36define <2 x bfloat> @cvt_rz_bf16x2_f32(float %f1, float %f2) {37; CHECK-LABEL: cvt_rz_bf16x2_f32(38; CHECK: {39; CHECK-NEXT: .reg .b32 %r<4>;40; CHECK-EMPTY:41; CHECK-NEXT: // %bb.0:42; CHECK-NEXT: ld.param.b32 %r1, [cvt_rz_bf16x2_f32_param_0];43; CHECK-NEXT: ld.param.b32 %r2, [cvt_rz_bf16x2_f32_param_1];44; CHECK-NEXT: cvt.rz.bf16x2.f32 %r3, %r1, %r2;45; CHECK-NEXT: st.param.b32 [func_retval0], %r3;46; CHECK-NEXT: ret;47 %val = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rz(float %f1, float %f2)48 ret <2 x bfloat> %val49}50 51define <2 x bfloat> @cvt_rz_relu_bf16x2_f32(float %f1, float %f2) {52; CHECK-LABEL: cvt_rz_relu_bf16x2_f32(53; CHECK: {54; CHECK-NEXT: .reg .b32 %r<4>;55; CHECK-EMPTY:56; CHECK-NEXT: // %bb.0:57; CHECK-NEXT: ld.param.b32 %r1, [cvt_rz_relu_bf16x2_f32_param_0];58; CHECK-NEXT: ld.param.b32 %r2, [cvt_rz_relu_bf16x2_f32_param_1];59; CHECK-NEXT: cvt.rz.relu.bf16x2.f32 %r3, %r1, %r2;60; CHECK-NEXT: st.param.b32 [func_retval0], %r3;61; CHECK-NEXT: ret;62 %val = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rz.relu(float %f1, float %f2)63 ret <2 x bfloat> %val64}65 66declare <2 x bfloat> @llvm.nvvm.ff2bf16x2.rn(float, float)67declare <2 x bfloat> @llvm.nvvm.ff2bf16x2.rn.relu(float, float)68declare <2 x bfloat> @llvm.nvvm.ff2bf16x2.rz(float, float)69declare <2 x bfloat> @llvm.nvvm.ff2bf16x2.rz.relu(float, float)70 71define <2 x half> @cvt_rn_f16x2_f32(float %f1, float %f2) {72; CHECK-LABEL: cvt_rn_f16x2_f32(73; CHECK: {74; CHECK-NEXT: .reg .b32 %r<4>;75; CHECK-EMPTY:76; CHECK-NEXT: // %bb.0:77; CHECK-NEXT: ld.param.b32 %r1, [cvt_rn_f16x2_f32_param_0];78; CHECK-NEXT: ld.param.b32 %r2, [cvt_rn_f16x2_f32_param_1];79; CHECK-NEXT: cvt.rn.f16x2.f32 %r3, %r1, %r2;80; CHECK-NEXT: st.param.b32 [func_retval0], %r3;81; CHECK-NEXT: ret;82 %val = call <2 x half> @llvm.nvvm.ff2f16x2.rn(float %f1, float %f2)83 ret <2 x half> %val84}85 86define <2 x half> @cvt_rn_relu_f16x2_f32(float %f1, float %f2) {87; CHECK-LABEL: cvt_rn_relu_f16x2_f32(88; CHECK: {89; CHECK-NEXT: .reg .b32 %r<4>;90; CHECK-EMPTY:91; CHECK-NEXT: // %bb.0:92; CHECK-NEXT: ld.param.b32 %r1, [cvt_rn_relu_f16x2_f32_param_0];93; CHECK-NEXT: ld.param.b32 %r2, [cvt_rn_relu_f16x2_f32_param_1];94; CHECK-NEXT: cvt.rn.relu.f16x2.f32 %r3, %r1, %r2;95; CHECK-NEXT: st.param.b32 [func_retval0], %r3;96; CHECK-NEXT: ret;97 %val = call <2 x half> @llvm.nvvm.ff2f16x2.rn.relu(float %f1, float %f2)98 ret <2 x half> %val99}100 101define <2 x half> @cvt_rz_f16x2_f32(float %f1, float %f2) {102; CHECK-LABEL: cvt_rz_f16x2_f32(103; CHECK: {104; CHECK-NEXT: .reg .b32 %r<4>;105; CHECK-EMPTY:106; CHECK-NEXT: // %bb.0:107; CHECK-NEXT: ld.param.b32 %r1, [cvt_rz_f16x2_f32_param_0];108; CHECK-NEXT: ld.param.b32 %r2, [cvt_rz_f16x2_f32_param_1];109; CHECK-NEXT: cvt.rz.f16x2.f32 %r3, %r1, %r2;110; CHECK-NEXT: st.param.b32 [func_retval0], %r3;111; CHECK-NEXT: ret;112 %val = call <2 x half> @llvm.nvvm.ff2f16x2.rz(float %f1, float %f2)113 ret <2 x half> %val114}115 116define <2 x half> @cvt_rz_relu_f16x2_f32(float %f1, float %f2) {117; CHECK-LABEL: cvt_rz_relu_f16x2_f32(118; CHECK: {119; CHECK-NEXT: .reg .b32 %r<4>;120; CHECK-EMPTY:121; CHECK-NEXT: // %bb.0:122; CHECK-NEXT: ld.param.b32 %r1, [cvt_rz_relu_f16x2_f32_param_0];123; CHECK-NEXT: ld.param.b32 %r2, [cvt_rz_relu_f16x2_f32_param_1];124; CHECK-NEXT: cvt.rz.relu.f16x2.f32 %r3, %r1, %r2;125; CHECK-NEXT: st.param.b32 [func_retval0], %r3;126; CHECK-NEXT: ret;127 %val = call <2 x half> @llvm.nvvm.ff2f16x2.rz.relu(float %f1, float %f2)128 ret <2 x half> %val129}130 131declare <2 x half> @llvm.nvvm.ff2f16x2.rn(float, float)132declare <2 x half> @llvm.nvvm.ff2f16x2.rn.relu(float, float)133declare <2 x half> @llvm.nvvm.ff2f16x2.rz(float, float)134declare <2 x half> @llvm.nvvm.ff2f16x2.rz.relu(float, float)135 136define bfloat @cvt_rn_bf16_f32(float %f1) {137; CHECK-LABEL: cvt_rn_bf16_f32(138; CHECK: {139; CHECK-NEXT: .reg .b16 %rs<2>;140; CHECK-NEXT: .reg .b32 %r<2>;141; CHECK-EMPTY:142; CHECK-NEXT: // %bb.0:143; CHECK-NEXT: ld.param.b32 %r1, [cvt_rn_bf16_f32_param_0];144; CHECK-NEXT: cvt.rn.bf16.f32 %rs1, %r1;145; CHECK-NEXT: st.param.b16 [func_retval0], %rs1;146; CHECK-NEXT: ret;147 %val = call bfloat @llvm.nvvm.f2bf16.rn(float %f1)148 ret bfloat %val149}150 151define bfloat @cvt_rn_relu_bf16_f32(float %f1) {152; CHECK-LABEL: cvt_rn_relu_bf16_f32(153; CHECK: {154; CHECK-NEXT: .reg .b16 %rs<2>;155; CHECK-NEXT: .reg .b32 %r<2>;156; CHECK-EMPTY:157; CHECK-NEXT: // %bb.0:158; CHECK-NEXT: ld.param.b32 %r1, [cvt_rn_relu_bf16_f32_param_0];159; CHECK-NEXT: cvt.rn.relu.bf16.f32 %rs1, %r1;160; CHECK-NEXT: st.param.b16 [func_retval0], %rs1;161; CHECK-NEXT: ret;162 %val = call bfloat @llvm.nvvm.f2bf16.rn.relu(float %f1)163 ret bfloat %val164}165 166define bfloat @cvt_rz_bf16_f32(float %f1) {167; CHECK-LABEL: cvt_rz_bf16_f32(168; CHECK: {169; CHECK-NEXT: .reg .b16 %rs<2>;170; CHECK-NEXT: .reg .b32 %r<2>;171; CHECK-EMPTY:172; CHECK-NEXT: // %bb.0:173; CHECK-NEXT: ld.param.b32 %r1, [cvt_rz_bf16_f32_param_0];174; CHECK-NEXT: cvt.rz.bf16.f32 %rs1, %r1;175; CHECK-NEXT: st.param.b16 [func_retval0], %rs1;176; CHECK-NEXT: ret;177 %val = call bfloat @llvm.nvvm.f2bf16.rz(float %f1)178 ret bfloat %val179}180 181define bfloat @cvt_rz_relu_bf16_f32(float %f1) {182; CHECK-LABEL: cvt_rz_relu_bf16_f32(183; CHECK: {184; CHECK-NEXT: .reg .b16 %rs<2>;185; CHECK-NEXT: .reg .b32 %r<2>;186; CHECK-EMPTY:187; CHECK-NEXT: // %bb.0:188; CHECK-NEXT: ld.param.b32 %r1, [cvt_rz_relu_bf16_f32_param_0];189; CHECK-NEXT: cvt.rz.relu.bf16.f32 %rs1, %r1;190; CHECK-NEXT: st.param.b16 [func_retval0], %rs1;191; CHECK-NEXT: ret;192 %val = call bfloat @llvm.nvvm.f2bf16.rz.relu(float %f1)193 ret bfloat %val194}195 196declare bfloat @llvm.nvvm.f2bf16.rn(float)197declare bfloat @llvm.nvvm.f2bf16.rn.relu(float)198declare bfloat @llvm.nvvm.f2bf16.rz(float)199declare bfloat @llvm.nvvm.f2bf16.rz.relu(float)200 201define half @cvt_rn_f16_f32(float %f1) {202; CHECK-LABEL: cvt_rn_f16_f32(203; CHECK: {204; CHECK-NEXT: .reg .b16 %rs<2>;205; CHECK-NEXT: .reg .b32 %r<2>;206; CHECK-EMPTY:207; CHECK-NEXT: // %bb.0:208; CHECK-NEXT: ld.param.b32 %r1, [cvt_rn_f16_f32_param_0];209; CHECK-NEXT: cvt.rn.f16.f32 %rs1, %r1;210; CHECK-NEXT: st.param.b16 [func_retval0], %rs1;211; CHECK-NEXT: ret;212 %val = call half @llvm.nvvm.f2f16.rn(float %f1)213 ret half %val214}215 216define half @cvt_rn_relu_f16_f32(float %f1) {217; CHECK-LABEL: cvt_rn_relu_f16_f32(218; CHECK: {219; CHECK-NEXT: .reg .b16 %rs<2>;220; CHECK-NEXT: .reg .b32 %r<2>;221; CHECK-EMPTY:222; CHECK-NEXT: // %bb.0:223; CHECK-NEXT: ld.param.b32 %r1, [cvt_rn_relu_f16_f32_param_0];224; CHECK-NEXT: cvt.rn.relu.f16.f32 %rs1, %r1;225; CHECK-NEXT: st.param.b16 [func_retval0], %rs1;226; CHECK-NEXT: ret;227 %val = call half @llvm.nvvm.f2f16.rn.relu(float %f1)228 ret half %val229}230 231define half @cvt_rz_f16_f32(float %f1) {232; CHECK-LABEL: cvt_rz_f16_f32(233; CHECK: {234; CHECK-NEXT: .reg .b16 %rs<2>;235; CHECK-NEXT: .reg .b32 %r<2>;236; CHECK-EMPTY:237; CHECK-NEXT: // %bb.0:238; CHECK-NEXT: ld.param.b32 %r1, [cvt_rz_f16_f32_param_0];239; CHECK-NEXT: cvt.rz.f16.f32 %rs1, %r1;240; CHECK-NEXT: st.param.b16 [func_retval0], %rs1;241; CHECK-NEXT: ret;242 %val = call half @llvm.nvvm.f2f16.rz(float %f1)243 ret half %val244}245 246define half @cvt_rz_relu_f16_f32(float %f1) {247; CHECK-LABEL: cvt_rz_relu_f16_f32(248; CHECK: {249; CHECK-NEXT: .reg .b16 %rs<2>;250; CHECK-NEXT: .reg .b32 %r<2>;251; CHECK-EMPTY:252; CHECK-NEXT: // %bb.0:253; CHECK-NEXT: ld.param.b32 %r1, [cvt_rz_relu_f16_f32_param_0];254; CHECK-NEXT: cvt.rz.relu.f16.f32 %rs1, %r1;255; CHECK-NEXT: st.param.b16 [func_retval0], %rs1;256; CHECK-NEXT: ret;257 %val = call half @llvm.nvvm.f2f16.rz.relu(float %f1)258 ret half %val259}260 261declare half @llvm.nvvm.f2f16.rn(float)262declare half @llvm.nvvm.f2f16.rn.relu(float)263declare half @llvm.nvvm.f2f16.rz(float)264declare half @llvm.nvvm.f2f16.rz.relu(float)265 266define i32 @cvt_rna_tf32_f32(float %f1) {267; CHECK-LABEL: cvt_rna_tf32_f32(268; CHECK: {269; CHECK-NEXT: .reg .b32 %r<3>;270; CHECK-EMPTY:271; CHECK-NEXT: // %bb.0:272; CHECK-NEXT: ld.param.b32 %r1, [cvt_rna_tf32_f32_param_0];273; CHECK-NEXT: cvt.rna.tf32.f32 %r2, %r1;274; CHECK-NEXT: st.param.b32 [func_retval0], %r2;275; CHECK-NEXT: ret;276 %val = call i32 @llvm.nvvm.f2tf32.rna(float %f1)277 ret i32 %val278}279 280declare i32 @llvm.nvvm.f2tf32.rna(float)281 282 283define <2 x bfloat> @fold_ff2bf16x2(float %lo, float %hi) {284; CHECK-LABEL: fold_ff2bf16x2(285; CHECK: {286; CHECK-NEXT: .reg .b32 %r<4>;287; CHECK-EMPTY:288; CHECK-NEXT: // %bb.0:289; CHECK-NEXT: ld.param.b32 %r1, [fold_ff2bf16x2_param_0];290; CHECK-NEXT: ld.param.b32 %r2, [fold_ff2bf16x2_param_1];291; CHECK-NEXT: cvt.rn.bf16x2.f32 %r3, %r2, %r1;292; CHECK-NEXT: st.param.b32 [func_retval0], %r3;293; CHECK-NEXT: ret;294 %loh = fptrunc float %lo to bfloat295 %hih = fptrunc float %hi to bfloat296 %v0 = insertelement <2 x bfloat> poison, bfloat %loh, i64 0297 %v1 = insertelement <2 x bfloat> %v0, bfloat %hih, i64 1298 ret <2 x bfloat> %v1299}300 301define <2 x half> @fold_ff2f16x2(float %lo, float %hi) {302; CHECK-LABEL: fold_ff2f16x2(303; CHECK: {304; CHECK-NEXT: .reg .b32 %r<4>;305; CHECK-EMPTY:306; CHECK-NEXT: // %bb.0:307; CHECK-NEXT: ld.param.b32 %r1, [fold_ff2f16x2_param_0];308; CHECK-NEXT: ld.param.b32 %r2, [fold_ff2f16x2_param_1];309; CHECK-NEXT: cvt.rn.f16x2.f32 %r3, %r2, %r1;310; CHECK-NEXT: st.param.b32 [func_retval0], %r3;311; CHECK-NEXT: ret;312 %loh = fptrunc float %lo to half313 %hih = fptrunc float %hi to half314 %v0 = insertelement <2 x half> poison, half %loh, i64 0315 %v1 = insertelement <2 x half> %v0, half %hih, i64 1316 ret <2 x half> %v1317}318