363 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_100a -mattr=+ptx86 | FileCheck %s3; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_101a -mattr=+ptx86 | FileCheck %s4; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_120a -mattr=+ptx86 | FileCheck %s5; RUN: %if ptxas-sm_100a && ptxas-isa-8.6 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100a -mattr=+ptx86 | %ptxas-verify -arch=sm_100a %}6; RUN: %if ptxas-sm_101a && ptxas-isa-8.6 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_101a -mattr=+ptx86 | %ptxas-verify -arch=sm_101a %}7; RUN: %if ptxas-sm_120a && ptxas-isa-8.6 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_120a -mattr=+ptx86 | %ptxas-verify -arch=sm_120a %}8 9define i16 @cvt_rn_sf_e2m3x2_f32(float %f1, float %f2) {10; CHECK-LABEL: cvt_rn_sf_e2m3x2_f32(11; CHECK: {12; CHECK-NEXT: .reg .b16 %rs<2>;13; CHECK-NEXT: .reg .b32 %r<4>;14; CHECK-EMPTY:15; CHECK-NEXT: // %bb.0:16; CHECK-NEXT: ld.param.b32 %r1, [cvt_rn_sf_e2m3x2_f32_param_0];17; CHECK-NEXT: ld.param.b32 %r2, [cvt_rn_sf_e2m3x2_f32_param_1];18; CHECK-NEXT: cvt.rn.satfinite.e2m3x2.f32 %rs1, %r1, %r2;19; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;20; CHECK-NEXT: st.param.b32 [func_retval0], %r3;21; CHECK-NEXT: ret;22 %val = call i16 @llvm.nvvm.ff.to.e2m3x2.rn.satfinite(float %f1, float %f2)23 ret i16 %val24}25 26define i16 @cvt_rn_relu_sf_e2m3x2_f32(float %f1, float %f2) {27; CHECK-LABEL: cvt_rn_relu_sf_e2m3x2_f32(28; CHECK: {29; CHECK-NEXT: .reg .b16 %rs<2>;30; CHECK-NEXT: .reg .b32 %r<4>;31; CHECK-EMPTY:32; CHECK-NEXT: // %bb.0:33; CHECK-NEXT: ld.param.b32 %r1, [cvt_rn_relu_sf_e2m3x2_f32_param_0];34; CHECK-NEXT: ld.param.b32 %r2, [cvt_rn_relu_sf_e2m3x2_f32_param_1];35; CHECK-NEXT: cvt.rn.satfinite.relu.e2m3x2.f32 %rs1, %r1, %r2;36; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;37; CHECK-NEXT: st.param.b32 [func_retval0], %r3;38; CHECK-NEXT: ret;39 %val = call i16 @llvm.nvvm.ff.to.e2m3x2.rn.relu.satfinite(float %f1, float %f2)40 ret i16 %val41}42 43define i16 @cvt_rn_sf_e3m2x2_f32(float %f1, float %f2) {44; CHECK-LABEL: cvt_rn_sf_e3m2x2_f32(45; CHECK: {46; CHECK-NEXT: .reg .b16 %rs<2>;47; CHECK-NEXT: .reg .b32 %r<4>;48; CHECK-EMPTY:49; CHECK-NEXT: // %bb.0:50; CHECK-NEXT: ld.param.b32 %r1, [cvt_rn_sf_e3m2x2_f32_param_0];51; CHECK-NEXT: ld.param.b32 %r2, [cvt_rn_sf_e3m2x2_f32_param_1];52; CHECK-NEXT: cvt.rn.satfinite.e3m2x2.f32 %rs1, %r1, %r2;53; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;54; CHECK-NEXT: st.param.b32 [func_retval0], %r3;55; CHECK-NEXT: ret;56 %val = call i16 @llvm.nvvm.ff.to.e3m2x2.rn.satfinite(float %f1, float %f2)57 ret i16 %val58}59 60define i16 @cvt_rn_relu_sf_e3m2x2_f32(float %f1, float %f2) {61; CHECK-LABEL: cvt_rn_relu_sf_e3m2x2_f32(62; CHECK: {63; CHECK-NEXT: .reg .b16 %rs<2>;64; CHECK-NEXT: .reg .b32 %r<4>;65; CHECK-EMPTY:66; CHECK-NEXT: // %bb.0:67; CHECK-NEXT: ld.param.b32 %r1, [cvt_rn_relu_sf_e3m2x2_f32_param_0];68; CHECK-NEXT: ld.param.b32 %r2, [cvt_rn_relu_sf_e3m2x2_f32_param_1];69; CHECK-NEXT: cvt.rn.satfinite.relu.e3m2x2.f32 %rs1, %r1, %r2;70; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;71; CHECK-NEXT: st.param.b32 [func_retval0], %r3;72; CHECK-NEXT: ret;73 %val = call i16 @llvm.nvvm.ff.to.e3m2x2.rn.relu.satfinite(float %f1, float %f2)74 ret i16 %val75}76 77define <2 x half> @cvt_rn_f16x2_e2m3x2(i16 %in) {78; CHECK-LABEL: cvt_rn_f16x2_e2m3x2(79; CHECK: {80; CHECK-NEXT: .reg .b16 %rs<2>;81; CHECK-NEXT: .reg .b32 %r<2>;82; CHECK-EMPTY:83; CHECK-NEXT: // %bb.0:84; CHECK-NEXT: ld.param.b16 %rs1, [cvt_rn_f16x2_e2m3x2_param_0];85; CHECK-NEXT: cvt.rn.f16x2.e2m3x2 %r1, %rs1;86; CHECK-NEXT: st.param.b32 [func_retval0], %r1;87; CHECK-NEXT: ret;88 %val = call <2 x half> @llvm.nvvm.e2m3x2.to.f16x2.rn(i16 %in)89 ret <2 x half> %val90}91 92define <2 x half> @cvt_rn_relu_f16x2_e2m3x2_relu(i16 %in) {93; CHECK-LABEL: cvt_rn_relu_f16x2_e2m3x2_relu(94; CHECK: {95; CHECK-NEXT: .reg .b16 %rs<2>;96; CHECK-NEXT: .reg .b32 %r<2>;97; CHECK-EMPTY:98; CHECK-NEXT: // %bb.0:99; CHECK-NEXT: ld.param.b16 %rs1, [cvt_rn_relu_f16x2_e2m3x2_relu_param_0];100; CHECK-NEXT: cvt.rn.relu.f16x2.e2m3x2 %r1, %rs1;101; CHECK-NEXT: st.param.b32 [func_retval0], %r1;102; CHECK-NEXT: ret;103 %val = call <2 x half> @llvm.nvvm.e2m3x2.to.f16x2.rn.relu(i16 %in)104 ret <2 x half> %val105}106 107define <2 x half> @cvt_rn_f16x2_e3m2x2(i16 %in) {108; CHECK-LABEL: cvt_rn_f16x2_e3m2x2(109; CHECK: {110; CHECK-NEXT: .reg .b16 %rs<2>;111; CHECK-NEXT: .reg .b32 %r<2>;112; CHECK-EMPTY:113; CHECK-NEXT: // %bb.0:114; CHECK-NEXT: ld.param.b16 %rs1, [cvt_rn_f16x2_e3m2x2_param_0];115; CHECK-NEXT: cvt.rn.f16x2.e3m2x2 %r1, %rs1;116; CHECK-NEXT: st.param.b32 [func_retval0], %r1;117; CHECK-NEXT: ret;118 %val = call <2 x half> @llvm.nvvm.e3m2x2.to.f16x2.rn(i16 %in)119 ret <2 x half> %val120}121 122define <2 x half> @cvt_rn_relu_f16x2_e3m2x2(i16 %in) {123; CHECK-LABEL: cvt_rn_relu_f16x2_e3m2x2(124; CHECK: {125; CHECK-NEXT: .reg .b16 %rs<2>;126; CHECK-NEXT: .reg .b32 %r<2>;127; CHECK-EMPTY:128; CHECK-NEXT: // %bb.0:129; CHECK-NEXT: ld.param.b16 %rs1, [cvt_rn_relu_f16x2_e3m2x2_param_0];130; CHECK-NEXT: cvt.rn.relu.f16x2.e3m2x2 %r1, %rs1;131; CHECK-NEXT: st.param.b32 [func_retval0], %r1;132; CHECK-NEXT: ret;133 %val = call <2 x half> @llvm.nvvm.e3m2x2.to.f16x2.rn.relu(i16 %in)134 ret <2 x half> %val135}136 137define i16 @cvt_rz_ue8m0x2_f32(float %f1, float %f2) {138; CHECK-LABEL: cvt_rz_ue8m0x2_f32(139; CHECK: {140; CHECK-NEXT: .reg .b16 %rs<2>;141; CHECK-NEXT: .reg .b32 %r<4>;142; CHECK-EMPTY:143; CHECK-NEXT: // %bb.0:144; CHECK-NEXT: ld.param.b32 %r1, [cvt_rz_ue8m0x2_f32_param_0];145; CHECK-NEXT: ld.param.b32 %r2, [cvt_rz_ue8m0x2_f32_param_1];146; CHECK-NEXT: cvt.rz.ue8m0x2.f32 %rs1, %r1, %r2;147; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;148; CHECK-NEXT: st.param.b32 [func_retval0], %r3;149; CHECK-NEXT: ret;150 %val = call i16 @llvm.nvvm.ff.to.ue8m0x2.rz(float %f1, float %f2)151 ret i16 %val152}153 154define i16 @cvt_rz_sf_ue8m0x2_f32(float %f1, float %f2) {155; CHECK-LABEL: cvt_rz_sf_ue8m0x2_f32(156; CHECK: {157; CHECK-NEXT: .reg .b16 %rs<2>;158; CHECK-NEXT: .reg .b32 %r<4>;159; CHECK-EMPTY:160; CHECK-NEXT: // %bb.0:161; CHECK-NEXT: ld.param.b32 %r1, [cvt_rz_sf_ue8m0x2_f32_param_0];162; CHECK-NEXT: ld.param.b32 %r2, [cvt_rz_sf_ue8m0x2_f32_param_1];163; CHECK-NEXT: cvt.rz.satfinite.ue8m0x2.f32 %rs1, %r1, %r2;164; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;165; CHECK-NEXT: st.param.b32 [func_retval0], %r3;166; CHECK-NEXT: ret;167 %val = call i16 @llvm.nvvm.ff.to.ue8m0x2.rz.satfinite(float %f1, float %f2)168 ret i16 %val169}170 171define i16 @cvt_rp_ue8m0x2_f32(float %f1, float %f2) {172; CHECK-LABEL: cvt_rp_ue8m0x2_f32(173; CHECK: {174; CHECK-NEXT: .reg .b16 %rs<2>;175; CHECK-NEXT: .reg .b32 %r<4>;176; CHECK-EMPTY:177; CHECK-NEXT: // %bb.0:178; CHECK-NEXT: ld.param.b32 %r1, [cvt_rp_ue8m0x2_f32_param_0];179; CHECK-NEXT: ld.param.b32 %r2, [cvt_rp_ue8m0x2_f32_param_1];180; CHECK-NEXT: cvt.rp.ue8m0x2.f32 %rs1, %r1, %r2;181; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;182; CHECK-NEXT: st.param.b32 [func_retval0], %r3;183; CHECK-NEXT: ret;184 %val = call i16 @llvm.nvvm.ff.to.ue8m0x2.rp(float %f1, float %f2)185 ret i16 %val186}187 188define i16 @cvt_rp_sf_ue8m0x2_f32(float %f1, float %f2) {189; CHECK-LABEL: cvt_rp_sf_ue8m0x2_f32(190; CHECK: {191; CHECK-NEXT: .reg .b16 %rs<2>;192; CHECK-NEXT: .reg .b32 %r<4>;193; CHECK-EMPTY:194; CHECK-NEXT: // %bb.0:195; CHECK-NEXT: ld.param.b32 %r1, [cvt_rp_sf_ue8m0x2_f32_param_0];196; CHECK-NEXT: ld.param.b32 %r2, [cvt_rp_sf_ue8m0x2_f32_param_1];197; CHECK-NEXT: cvt.rp.satfinite.ue8m0x2.f32 %rs1, %r1, %r2;198; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;199; CHECK-NEXT: st.param.b32 [func_retval0], %r3;200; CHECK-NEXT: ret;201 %val = call i16 @llvm.nvvm.ff.to.ue8m0x2.rp.satfinite(float %f1, float %f2)202 ret i16 %val203}204 205define i16 @cvt_rz_ue8m0x2_bf16x2(<2 x bfloat> %in) {206; CHECK-LABEL: cvt_rz_ue8m0x2_bf16x2(207; CHECK: {208; CHECK-NEXT: .reg .b16 %rs<2>;209; CHECK-NEXT: .reg .b32 %r<3>;210; CHECK-EMPTY:211; CHECK-NEXT: // %bb.0:212; CHECK-NEXT: ld.param.b32 %r1, [cvt_rz_ue8m0x2_bf16x2_param_0];213; CHECK-NEXT: cvt.rz.ue8m0x2.bf16x2 %rs1, %r1;214; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;215; CHECK-NEXT: st.param.b32 [func_retval0], %r2;216; CHECK-NEXT: ret;217 %val = call i16 @llvm.nvvm.bf16x2.to.ue8m0x2.rz(<2 x bfloat> %in)218 ret i16 %val219}220 221define i16 @cvt_rz_sf_ue8m0x2_bf16x2(<2 x bfloat> %in) {222; CHECK-LABEL: cvt_rz_sf_ue8m0x2_bf16x2(223; CHECK: {224; CHECK-NEXT: .reg .b16 %rs<2>;225; CHECK-NEXT: .reg .b32 %r<3>;226; CHECK-EMPTY:227; CHECK-NEXT: // %bb.0:228; CHECK-NEXT: ld.param.b32 %r1, [cvt_rz_sf_ue8m0x2_bf16x2_param_0];229; CHECK-NEXT: cvt.rz.satfinite.ue8m0x2.bf16x2 %rs1, %r1;230; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;231; CHECK-NEXT: st.param.b32 [func_retval0], %r2;232; CHECK-NEXT: ret;233 %val = call i16 @llvm.nvvm.bf16x2.to.ue8m0x2.rz.satfinite(<2 x bfloat> %in)234 ret i16 %val235}236 237define i16 @cvt_rp_ue8m0x2_bf16x2(<2 x bfloat> %in) {238; CHECK-LABEL: cvt_rp_ue8m0x2_bf16x2(239; CHECK: {240; CHECK-NEXT: .reg .b16 %rs<2>;241; CHECK-NEXT: .reg .b32 %r<3>;242; CHECK-EMPTY:243; CHECK-NEXT: // %bb.0:244; CHECK-NEXT: ld.param.b32 %r1, [cvt_rp_ue8m0x2_bf16x2_param_0];245; CHECK-NEXT: cvt.rp.ue8m0x2.bf16x2 %rs1, %r1;246; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;247; CHECK-NEXT: st.param.b32 [func_retval0], %r2;248; CHECK-NEXT: ret;249 %val = call i16 @llvm.nvvm.bf16x2.to.ue8m0x2.rp(<2 x bfloat> %in)250 ret i16 %val251}252 253define i16 @cvt_rp_sf_ue8m0x2_bf16x2(<2 x bfloat> %in) {254; CHECK-LABEL: cvt_rp_sf_ue8m0x2_bf16x2(255; CHECK: {256; CHECK-NEXT: .reg .b16 %rs<2>;257; CHECK-NEXT: .reg .b32 %r<3>;258; CHECK-EMPTY:259; CHECK-NEXT: // %bb.0:260; CHECK-NEXT: ld.param.b32 %r1, [cvt_rp_sf_ue8m0x2_bf16x2_param_0];261; CHECK-NEXT: cvt.rp.satfinite.ue8m0x2.bf16x2 %rs1, %r1;262; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;263; CHECK-NEXT: st.param.b32 [func_retval0], %r2;264; CHECK-NEXT: ret;265 %val = call i16 @llvm.nvvm.bf16x2.to.ue8m0x2.rp.satfinite(<2 x bfloat> %in)266 ret i16 %val267}268 269define <2 x bfloat> @cvt_bf16x2_ue8m0x2(i16 %in) {270; CHECK-LABEL: cvt_bf16x2_ue8m0x2(271; CHECK: {272; CHECK-NEXT: .reg .b16 %rs<2>;273; CHECK-NEXT: .reg .b32 %r<2>;274; CHECK-EMPTY:275; CHECK-NEXT: // %bb.0:276; CHECK-NEXT: ld.param.b16 %rs1, [cvt_bf16x2_ue8m0x2_param_0];277; CHECK-NEXT: cvt.rn.bf16x2.ue8m0x2 %r1, %rs1;278; CHECK-NEXT: st.param.b32 [func_retval0], %r1;279; CHECK-NEXT: ret;280 %val = call <2 x bfloat> @llvm.nvvm.ue8m0x2.to.bf16x2(i16 %in)281 ret <2 x bfloat> %val282}283 284define i16 @cvt_rn_sf_e2m1x2_f32(float %f1, float %f2) {285; CHECK-LABEL: cvt_rn_sf_e2m1x2_f32(286; CHECK: {287; CHECK-NEXT: .reg .b16 %rs<2>;288; CHECK-NEXT: .reg .b32 %r<4>;289; CHECK-EMPTY:290; CHECK-NEXT: // %bb.0:291; CHECK-NEXT: ld.param.b32 %r1, [cvt_rn_sf_e2m1x2_f32_param_0];292; CHECK-NEXT: ld.param.b32 %r2, [cvt_rn_sf_e2m1x2_f32_param_1];293; CHECK-NEXT: {294; CHECK-NEXT: .reg .b8 %e2m1x2_out;295; CHECK-NEXT: cvt.rn.satfinite.e2m1x2.f32 %e2m1x2_out, %r1, %r2;296; CHECK-NEXT: cvt.u16.u8 %rs1, %e2m1x2_out;297; CHECK-NEXT: }298; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;299; CHECK-NEXT: st.param.b32 [func_retval0], %r3;300; CHECK-NEXT: ret;301 %val = call i16 @llvm.nvvm.ff.to.e2m1x2.rn.satfinite(float %f1, float %f2)302 ret i16 %val303}304 305define i16 @cvt_rn_relu_sf_e2m1x2_f32(float %f1, float %f2) {306; CHECK-LABEL: cvt_rn_relu_sf_e2m1x2_f32(307; CHECK: {308; CHECK-NEXT: .reg .b16 %rs<2>;309; CHECK-NEXT: .reg .b32 %r<4>;310; CHECK-EMPTY:311; CHECK-NEXT: // %bb.0:312; CHECK-NEXT: ld.param.b32 %r1, [cvt_rn_relu_sf_e2m1x2_f32_param_0];313; CHECK-NEXT: ld.param.b32 %r2, [cvt_rn_relu_sf_e2m1x2_f32_param_1];314; CHECK-NEXT: {315; CHECK-NEXT: .reg .b8 %e2m1x2_out;316; CHECK-NEXT: cvt.rn.satfinite.relu.e2m1x2.f32 %e2m1x2_out, %r1, %r2;317; CHECK-NEXT: cvt.u16.u8 %rs1, %e2m1x2_out;318; CHECK-NEXT: }319; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;320; CHECK-NEXT: st.param.b32 [func_retval0], %r3;321; CHECK-NEXT: ret;322 %val = call i16 @llvm.nvvm.ff.to.e2m1x2.rn.relu.satfinite(float %f1, float %f2)323 ret i16 %val324}325 326define <2 x half> @cvt_rn_f16x2_e2m1x2(i16 %in) {327; CHECK-LABEL: cvt_rn_f16x2_e2m1x2(328; CHECK: {329; CHECK-NEXT: .reg .b16 %rs<2>;330; CHECK-NEXT: .reg .b32 %r<2>;331; CHECK-EMPTY:332; CHECK-NEXT: // %bb.0:333; CHECK-NEXT: ld.param.b16 %rs1, [cvt_rn_f16x2_e2m1x2_param_0];334; CHECK-NEXT: {335; CHECK-NEXT: .reg .b8 %e2m1x2_in;336; CHECK-NEXT: cvt.u8.u16 %e2m1x2_in, %rs1;337; CHECK-NEXT: cvt.rn.f16x2.e2m1x2 %r1, %e2m1x2_in;338; CHECK-NEXT: }339; CHECK-NEXT: st.param.b32 [func_retval0], %r1;340; CHECK-NEXT: ret;341 %val = call <2 x half> @llvm.nvvm.e2m1x2.to.f16x2.rn(i16 %in)342 ret <2 x half> %val343}344 345define <2 x half> @cvt_rn_relu_f16x2_e2m1x2(i16 %in) {346; CHECK-LABEL: cvt_rn_relu_f16x2_e2m1x2(347; CHECK: {348; CHECK-NEXT: .reg .b16 %rs<2>;349; CHECK-NEXT: .reg .b32 %r<2>;350; CHECK-EMPTY:351; CHECK-NEXT: // %bb.0:352; CHECK-NEXT: ld.param.b16 %rs1, [cvt_rn_relu_f16x2_e2m1x2_param_0];353; CHECK-NEXT: {354; CHECK-NEXT: .reg .b8 %e2m1x2_in;355; CHECK-NEXT: cvt.u8.u16 %e2m1x2_in, %rs1;356; CHECK-NEXT: cvt.rn.relu.f16x2.e2m1x2 %r1, %e2m1x2_in;357; CHECK-NEXT: }358; CHECK-NEXT: st.param.b32 [func_retval0], %r1;359; CHECK-NEXT: ret;360 %val = call <2 x half> @llvm.nvvm.e2m1x2.to.f16x2.rn.relu(i16 %in)361 ret <2 x half> %val362}363