298 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_100a -mattr=+ptx87 | FileCheck %s3; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_103a -mattr=+ptx87 | FileCheck %s4; RUN: %if ptxas-sm_100a && ptxas-isa-8.7 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100a -mattr=+ptx87 | %ptxas-verify -arch=sm_100a %}5; RUN: %if ptxas-sm_103a && ptxas-isa-8.7 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_103a -mattr=+ptx87 | %ptxas-verify -arch=sm_103a %}6 7; F16X2 conversions8 9define <2 x half> @cvt_rs_f16x2_f32(float %f1, float %f2, i32 %rbits) {10; CHECK-LABEL: cvt_rs_f16x2_f32(11; CHECK: {12; CHECK-NEXT: .reg .b32 %r<5>;13; CHECK-EMPTY:14; CHECK-NEXT: // %bb.0:15; CHECK-NEXT: ld.param.b32 %r1, [cvt_rs_f16x2_f32_param_0];16; CHECK-NEXT: ld.param.b32 %r2, [cvt_rs_f16x2_f32_param_1];17; CHECK-NEXT: ld.param.b32 %r3, [cvt_rs_f16x2_f32_param_2];18; CHECK-NEXT: cvt.rs.f16x2.f32 %r4, %r1, %r2, %r3;19; CHECK-NEXT: st.param.b32 [func_retval0], %r4;20; CHECK-NEXT: ret;21 %val = call <2 x half> @llvm.nvvm.ff2f16x2.rs(float %f1, float %f2, i32 %rbits)22 ret <2 x half> %val23}24 25define <2 x half> @cvt_rs_relu_f16x2_f32(float %f1, float %f2, i32 %rbits) {26; CHECK-LABEL: cvt_rs_relu_f16x2_f32(27; CHECK: {28; CHECK-NEXT: .reg .b32 %r<5>;29; CHECK-EMPTY:30; CHECK-NEXT: // %bb.0:31; CHECK-NEXT: ld.param.b32 %r1, [cvt_rs_relu_f16x2_f32_param_0];32; CHECK-NEXT: ld.param.b32 %r2, [cvt_rs_relu_f16x2_f32_param_1];33; CHECK-NEXT: ld.param.b32 %r3, [cvt_rs_relu_f16x2_f32_param_2];34; CHECK-NEXT: cvt.rs.relu.f16x2.f32 %r4, %r1, %r2, %r3;35; CHECK-NEXT: st.param.b32 [func_retval0], %r4;36; CHECK-NEXT: ret;37 %val = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu(float %f1, float %f2, i32 %rbits)38 ret <2 x half> %val39}40 41define <2 x half> @cvt_rs_sf_f16x2_f32(float %f1, float %f2, i32 %rbits) {42; CHECK-LABEL: cvt_rs_sf_f16x2_f32(43; CHECK: {44; CHECK-NEXT: .reg .b32 %r<5>;45; CHECK-EMPTY:46; CHECK-NEXT: // %bb.0:47; CHECK-NEXT: ld.param.b32 %r1, [cvt_rs_sf_f16x2_f32_param_0];48; CHECK-NEXT: ld.param.b32 %r2, [cvt_rs_sf_f16x2_f32_param_1];49; CHECK-NEXT: ld.param.b32 %r3, [cvt_rs_sf_f16x2_f32_param_2];50; CHECK-NEXT: cvt.rs.satfinite.f16x2.f32 %r4, %r1, %r2, %r3;51; CHECK-NEXT: st.param.b32 [func_retval0], %r4;52; CHECK-NEXT: ret;53 %val = call <2 x half> @llvm.nvvm.ff2f16x2.rs.satfinite(float %f1, float %f2, i32 %rbits)54 ret <2 x half> %val55}56 57define <2 x half> @cvt_rs_relu_sf_f16x2_f32(float %f1, float %f2, i32 %rbits) {58; CHECK-LABEL: cvt_rs_relu_sf_f16x2_f32(59; CHECK: {60; CHECK-NEXT: .reg .b32 %r<5>;61; CHECK-EMPTY:62; CHECK-NEXT: // %bb.0:63; CHECK-NEXT: ld.param.b32 %r1, [cvt_rs_relu_sf_f16x2_f32_param_0];64; CHECK-NEXT: ld.param.b32 %r2, [cvt_rs_relu_sf_f16x2_f32_param_1];65; CHECK-NEXT: ld.param.b32 %r3, [cvt_rs_relu_sf_f16x2_f32_param_2];66; CHECK-NEXT: cvt.rs.relu.satfinite.f16x2.f32 %r4, %r1, %r2, %r3;67; CHECK-NEXT: st.param.b32 [func_retval0], %r4;68; CHECK-NEXT: ret;69 %val = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu.satfinite(float %f1, float %f2, i32 %rbits)70 ret <2 x half> %val71}72 73; BF16X2 conversions74 75define <2 x bfloat> @cvt_rs_bf16x2_f32(float %f1, float %f2, i32 %rbits) {76; CHECK-LABEL: cvt_rs_bf16x2_f32(77; CHECK: {78; CHECK-NEXT: .reg .b32 %r<5>;79; CHECK-EMPTY:80; CHECK-NEXT: // %bb.0:81; CHECK-NEXT: ld.param.b32 %r1, [cvt_rs_bf16x2_f32_param_0];82; CHECK-NEXT: ld.param.b32 %r2, [cvt_rs_bf16x2_f32_param_1];83; CHECK-NEXT: ld.param.b32 %r3, [cvt_rs_bf16x2_f32_param_2];84; CHECK-NEXT: cvt.rs.bf16x2.f32 %r4, %r1, %r2, %r3;85; CHECK-NEXT: st.param.b32 [func_retval0], %r4;86; CHECK-NEXT: ret;87 %val = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs(float %f1, float %f2, i32 %rbits)88 ret <2 x bfloat> %val89}90 91define <2 x bfloat> @cvt_rs_relu_bf16x2_f32(float %f1, float %f2, i32 %rbits) {92; CHECK-LABEL: cvt_rs_relu_bf16x2_f32(93; CHECK: {94; CHECK-NEXT: .reg .b32 %r<5>;95; CHECK-EMPTY:96; CHECK-NEXT: // %bb.0:97; CHECK-NEXT: ld.param.b32 %r1, [cvt_rs_relu_bf16x2_f32_param_0];98; CHECK-NEXT: ld.param.b32 %r2, [cvt_rs_relu_bf16x2_f32_param_1];99; CHECK-NEXT: ld.param.b32 %r3, [cvt_rs_relu_bf16x2_f32_param_2];100; CHECK-NEXT: cvt.rs.relu.bf16x2.f32 %r4, %r1, %r2, %r3;101; CHECK-NEXT: st.param.b32 [func_retval0], %r4;102; CHECK-NEXT: ret;103 %val = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu(float %f1, float %f2, i32 %rbits)104 ret <2 x bfloat> %val105}106 107define <2 x bfloat> @cvt_rs_sf_bf16x2_f32(float %f1, float %f2, i32 %rbits) {108; CHECK-LABEL: cvt_rs_sf_bf16x2_f32(109; CHECK: {110; CHECK-NEXT: .reg .b32 %r<5>;111; CHECK-EMPTY:112; CHECK-NEXT: // %bb.0:113; CHECK-NEXT: ld.param.b32 %r1, [cvt_rs_sf_bf16x2_f32_param_0];114; CHECK-NEXT: ld.param.b32 %r2, [cvt_rs_sf_bf16x2_f32_param_1];115; CHECK-NEXT: ld.param.b32 %r3, [cvt_rs_sf_bf16x2_f32_param_2];116; CHECK-NEXT: cvt.rs.satfinite.bf16x2.f32 %r4, %r1, %r2, %r3;117; CHECK-NEXT: st.param.b32 [func_retval0], %r4;118; CHECK-NEXT: ret;119 %val = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.satfinite(float %f1, float %f2, i32 %rbits)120 ret <2 x bfloat> %val121}122 123define <2 x bfloat> @cvt_rs_relu_sf_bf16x2_f32(float %f1, float %f2, i32 %rbits) {124; CHECK-LABEL: cvt_rs_relu_sf_bf16x2_f32(125; CHECK: {126; CHECK-NEXT: .reg .b32 %r<5>;127; CHECK-EMPTY:128; CHECK-NEXT: // %bb.0:129; CHECK-NEXT: ld.param.b32 %r1, [cvt_rs_relu_sf_bf16x2_f32_param_0];130; CHECK-NEXT: ld.param.b32 %r2, [cvt_rs_relu_sf_bf16x2_f32_param_1];131; CHECK-NEXT: ld.param.b32 %r3, [cvt_rs_relu_sf_bf16x2_f32_param_2];132; CHECK-NEXT: cvt.rs.relu.satfinite.bf16x2.f32 %r4, %r1, %r2, %r3;133; CHECK-NEXT: st.param.b32 [func_retval0], %r4;134; CHECK-NEXT: ret;135 %val = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu.satfinite(float %f1, float %f2, i32 %rbits)136 ret <2 x bfloat> %val137}138 139; F8X4 conversions140 141define <4 x i8> @cvt_rs_sf_e4m3x4_f32(<4 x float> %fvec, i32 %rbits) {142; CHECK-LABEL: cvt_rs_sf_e4m3x4_f32(143; CHECK: {144; CHECK-NEXT: .reg .b32 %r<7>;145; CHECK-EMPTY:146; CHECK-NEXT: // %bb.0:147; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [cvt_rs_sf_e4m3x4_f32_param_0];148; CHECK-NEXT: ld.param.b32 %r5, [cvt_rs_sf_e4m3x4_f32_param_1];149; CHECK-NEXT: cvt.rs.satfinite.e4m3x4.f32 %r6, {%r1, %r2, %r3, %r4}, %r5;150; CHECK-NEXT: st.param.b32 [func_retval0], %r6;151; CHECK-NEXT: ret;152 %val = call <4 x i8> @llvm.nvvm.f32x4.to.e4m3x4.rs.satfinite(<4 x float> %fvec, i32 %rbits)153 ret <4 x i8> %val154}155 156define <4 x i8> @cvt_rs_relu_sf_e4m3x4_f32(<4 x float> %fvec, i32 %rbits) {157; CHECK-LABEL: cvt_rs_relu_sf_e4m3x4_f32(158; CHECK: {159; CHECK-NEXT: .reg .b32 %r<7>;160; CHECK-EMPTY:161; CHECK-NEXT: // %bb.0:162; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [cvt_rs_relu_sf_e4m3x4_f32_param_0];163; CHECK-NEXT: ld.param.b32 %r5, [cvt_rs_relu_sf_e4m3x4_f32_param_1];164; CHECK-NEXT: cvt.rs.relu.satfinite.e4m3x4.f32 %r6, {%r1, %r2, %r3, %r4}, %r5;165; CHECK-NEXT: st.param.b32 [func_retval0], %r6;166; CHECK-NEXT: ret;167 %val = call <4 x i8> @llvm.nvvm.f32x4.to.e4m3x4.rs.relu.satfinite(<4 x float> %fvec, i32 %rbits)168 ret <4 x i8> %val169}170 171define <4 x i8> @cvt_rs_sf_e5m2x4_f32(<4 x float> %fvec, i32 %rbits) {172; CHECK-LABEL: cvt_rs_sf_e5m2x4_f32(173; CHECK: {174; CHECK-NEXT: .reg .b32 %r<7>;175; CHECK-EMPTY:176; CHECK-NEXT: // %bb.0:177; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [cvt_rs_sf_e5m2x4_f32_param_0];178; CHECK-NEXT: ld.param.b32 %r5, [cvt_rs_sf_e5m2x4_f32_param_1];179; CHECK-NEXT: cvt.rs.satfinite.e5m2x4.f32 %r6, {%r1, %r2, %r3, %r4}, %r5;180; CHECK-NEXT: st.param.b32 [func_retval0], %r6;181; CHECK-NEXT: ret;182 %val = call <4 x i8> @llvm.nvvm.f32x4.to.e5m2x4.rs.satfinite(<4 x float> %fvec, i32 %rbits)183 ret <4 x i8> %val184}185 186define <4 x i8> @cvt_rs_relu_sf_e5m2x4_f32(<4 x float> %fvec, i32 %rbits) {187; CHECK-LABEL: cvt_rs_relu_sf_e5m2x4_f32(188; CHECK: {189; CHECK-NEXT: .reg .b32 %r<7>;190; CHECK-EMPTY:191; CHECK-NEXT: // %bb.0:192; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [cvt_rs_relu_sf_e5m2x4_f32_param_0];193; CHECK-NEXT: ld.param.b32 %r5, [cvt_rs_relu_sf_e5m2x4_f32_param_1];194; CHECK-NEXT: cvt.rs.relu.satfinite.e5m2x4.f32 %r6, {%r1, %r2, %r3, %r4}, %r5;195; CHECK-NEXT: st.param.b32 [func_retval0], %r6;196; CHECK-NEXT: ret;197 %val = call <4 x i8> @llvm.nvvm.f32x4.to.e5m2x4.rs.relu.satfinite(<4 x float> %fvec, i32 %rbits)198 ret <4 x i8> %val199}200 201; F6X4 conversions202 203define <4 x i8> @cvt_rs_sf_e2m3x4_f32(<4 x float> %fvec, i32 %rbits) {204; CHECK-LABEL: cvt_rs_sf_e2m3x4_f32(205; CHECK: {206; CHECK-NEXT: .reg .b32 %r<7>;207; CHECK-EMPTY:208; CHECK-NEXT: // %bb.0:209; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [cvt_rs_sf_e2m3x4_f32_param_0];210; CHECK-NEXT: ld.param.b32 %r5, [cvt_rs_sf_e2m3x4_f32_param_1];211; CHECK-NEXT: cvt.rs.satfinite.e2m3x4.f32 %r6, {%r1, %r2, %r3, %r4}, %r5;212; CHECK-NEXT: st.param.b32 [func_retval0], %r6;213; CHECK-NEXT: ret;214 %val = call <4 x i8> @llvm.nvvm.f32x4.to.e2m3x4.rs.satfinite(<4 x float> %fvec, i32 %rbits)215 ret <4 x i8> %val216}217 218define <4 x i8> @cvt_rs_relu_sf_e2m3x4_f32(<4 x float> %fvec, i32 %rbits) {219; CHECK-LABEL: cvt_rs_relu_sf_e2m3x4_f32(220; CHECK: {221; CHECK-NEXT: .reg .b32 %r<7>;222; CHECK-EMPTY:223; CHECK-NEXT: // %bb.0:224; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [cvt_rs_relu_sf_e2m3x4_f32_param_0];225; CHECK-NEXT: ld.param.b32 %r5, [cvt_rs_relu_sf_e2m3x4_f32_param_1];226; CHECK-NEXT: cvt.rs.relu.satfinite.e2m3x4.f32 %r6, {%r1, %r2, %r3, %r4}, %r5;227; CHECK-NEXT: st.param.b32 [func_retval0], %r6;228; CHECK-NEXT: ret;229 %val = call <4 x i8> @llvm.nvvm.f32x4.to.e2m3x4.rs.relu.satfinite(<4 x float> %fvec, i32 %rbits)230 ret <4 x i8> %val231}232 233define <4 x i8> @cvt_rs_sf_e3m2x4_f32(<4 x float> %fvec, i32 %rbits) {234; CHECK-LABEL: cvt_rs_sf_e3m2x4_f32(235; CHECK: {236; CHECK-NEXT: .reg .b32 %r<7>;237; CHECK-EMPTY:238; CHECK-NEXT: // %bb.0:239; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [cvt_rs_sf_e3m2x4_f32_param_0];240; CHECK-NEXT: ld.param.b32 %r5, [cvt_rs_sf_e3m2x4_f32_param_1];241; CHECK-NEXT: cvt.rs.satfinite.e3m2x4.f32 %r6, {%r1, %r2, %r3, %r4}, %r5;242; CHECK-NEXT: st.param.b32 [func_retval0], %r6;243; CHECK-NEXT: ret;244 %val = call <4 x i8> @llvm.nvvm.f32x4.to.e3m2x4.rs.satfinite(<4 x float> %fvec, i32 %rbits)245 ret <4 x i8> %val246}247 248define <4 x i8> @cvt_rs_relu_sf_e3m2x4_f32(<4 x float> %fvec, i32 %rbits) {249; CHECK-LABEL: cvt_rs_relu_sf_e3m2x4_f32(250; CHECK: {251; CHECK-NEXT: .reg .b32 %r<7>;252; CHECK-EMPTY:253; CHECK-NEXT: // %bb.0:254; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [cvt_rs_relu_sf_e3m2x4_f32_param_0];255; CHECK-NEXT: ld.param.b32 %r5, [cvt_rs_relu_sf_e3m2x4_f32_param_1];256; CHECK-NEXT: cvt.rs.relu.satfinite.e3m2x4.f32 %r6, {%r1, %r2, %r3, %r4}, %r5;257; CHECK-NEXT: st.param.b32 [func_retval0], %r6;258; CHECK-NEXT: ret;259 %val = call <4 x i8> @llvm.nvvm.f32x4.to.e3m2x4.rs.relu.satfinite(<4 x float> %fvec, i32 %rbits)260 ret <4 x i8> %val261}262 263; F4X4 conversions264 265define i16 @cvt_rs_sf_e2m1x4_f32(<4 x float> %fvec, i32 %rbits) {266; CHECK-LABEL: cvt_rs_sf_e2m1x4_f32(267; CHECK: {268; CHECK-NEXT: .reg .b16 %rs<2>;269; CHECK-NEXT: .reg .b32 %r<7>;270; CHECK-EMPTY:271; CHECK-NEXT: // %bb.0:272; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [cvt_rs_sf_e2m1x4_f32_param_0];273; CHECK-NEXT: ld.param.b32 %r5, [cvt_rs_sf_e2m1x4_f32_param_1];274; CHECK-NEXT: cvt.rs.satfinite.e2m1x4.f32 %rs1, {%r1, %r2, %r3, %r4}, %r5;275; CHECK-NEXT: cvt.u32.u16 %r6, %rs1;276; CHECK-NEXT: st.param.b32 [func_retval0], %r6;277; CHECK-NEXT: ret;278 %val = call i16 @llvm.nvvm.f32x4.to.e2m1x4.rs.satfinite(<4 x float> %fvec, i32 %rbits)279 ret i16 %val280}281 282define i16 @cvt_rs_relu_sf_e2m1x4_f32(<4 x float> %fvec, i32 %rbits) {283; CHECK-LABEL: cvt_rs_relu_sf_e2m1x4_f32(284; CHECK: {285; CHECK-NEXT: .reg .b16 %rs<2>;286; CHECK-NEXT: .reg .b32 %r<7>;287; CHECK-EMPTY:288; CHECK-NEXT: // %bb.0:289; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [cvt_rs_relu_sf_e2m1x4_f32_param_0];290; CHECK-NEXT: ld.param.b32 %r5, [cvt_rs_relu_sf_e2m1x4_f32_param_1];291; CHECK-NEXT: cvt.rs.relu.satfinite.e2m1x4.f32 %rs1, {%r1, %r2, %r3, %r4}, %r5;292; CHECK-NEXT: cvt.u32.u16 %r6, %rs1;293; CHECK-NEXT: st.param.b32 [func_retval0], %r6;294; CHECK-NEXT: ret;295 %val = call i16 @llvm.nvvm.f32x4.to.e2m1x4.rs.relu.satfinite(<4 x float> %fvec, i32 %rbits)296 ret i16 %val297}298