Google Git
Sign in
llvm/llvm-project/e14b11b74dcafea1e2a71b9788b85b34ab4cf1e6/./llvm/test/CodeGen/NVPTX/arbitrary-fp-to-float.ll
blob: fd3d72119a03231c016b29b48eab81715ab9f169 [file]
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
; RUN: llc < %s -mtriple=nvptx64-unknown-unknown | FileCheck %s
; Test llvm.convert.from.arbitrary intrinsic expansion.
declare float @llvm.convert.from.arbitrary.fp.f32.i8(i8, metadata)
declare float @llvm.convert.from.arbitrary.fp.f32.i6(i6, metadata)
declare float @llvm.convert.from.arbitrary.fp.f32.i4(i4, metadata)
declare <2 x float> @llvm.convert.from.arbitrary.fp.v2f32.v2i8(<2 x i8>, metadata)
declare <3 x float> @llvm.convert.from.arbitrary.fp.v3f32.v3i8(<3 x i8>, metadata)
declare <4 x float> @llvm.convert.from.arbitrary.fp.v4f32.v4i4(<4 x i4>, metadata)
declare <4 x float> @llvm.convert.from.arbitrary.fp.v4f32.v4i8(<4 x i8>, metadata)
declare <2 x half> @llvm.convert.from.arbitrary.fp.v2f16.v2i8(<2 x i8>, metadata)
declare half @llvm.convert.from.arbitrary.fp.f16.i8(i8, metadata)
declare double @llvm.convert.from.arbitrary.fp.f64.i8(i8, metadata)
declare bfloat @llvm.convert.from.arbitrary.fp.bf16.i8(i8, metadata)
declare <2 x bfloat> @llvm.convert.from.arbitrary.fp.v2bf16.v2i8(<2 x i8>, metadata)
; Float8E5M2
; Layout: sign(1) exp(5) mant(2), bias=15
; Supports: Inf, NaN, signed zero, denormals
; Float8E5M2 normal: 0_01111_00 = 1.0
define float @from_f8e5m2_normal() {
; CHECK-LABEL: from_f8e5m2_normal(
; CHECK: {
; CHECK-EMPTY:
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: st.param.b32 [func_retval0], 1065353216;
; CHECK-NEXT: ret;
%r = call float @llvm.convert.from.arbitrary.fp.f32.i8(i8 60, metadata !"Float8E5M2")
ret float %r
}
; Float8E5M2 zero: 0_00000_00 = +0.0
define float @from_f8e5m2_zero() {
; CHECK-LABEL: from_f8e5m2_zero(
; CHECK: {
; CHECK-EMPTY:
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: st.param.b32 [func_retval0], 0;
; CHECK-NEXT: ret;
%r = call float @llvm.convert.from.arbitrary.fp.f32.i8(i8 0, metadata !"Float8E5M2")
ret float %r
}
; Float8E5M2 negative zero: 1_00000_00 = -0.0
define float @from_f8e5m2_neg_zero() {
; CHECK-LABEL: from_f8e5m2_neg_zero(
; CHECK: {
; CHECK-EMPTY:
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: st.param.b32 [func_retval0], -2147483648;
; CHECK-NEXT: ret;
%r = call float @llvm.convert.from.arbitrary.fp.f32.i8(i8 -128, metadata !"Float8E5M2")
ret float %r
}
; Float8E5M2 denorm: 0_00000_01 = 2^(-16)
define float @from_f8e5m2_denorm() {
; CHECK-LABEL: from_f8e5m2_denorm(
; CHECK: {
; CHECK-EMPTY:
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: st.param.b32 [func_retval0], 931135488;
; CHECK-NEXT: ret;
%r = call float @llvm.convert.from.arbitrary.fp.f32.i8(i8 1, metadata !"Float8E5M2")
ret float %r
}
; Float8E5M2 +Inf: 0_11111_00
define float @from_f8e5m2_inf() {
; CHECK-LABEL: from_f8e5m2_inf(
; CHECK: {
; CHECK-EMPTY:
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: st.param.b32 [func_retval0], 2139095040;
; CHECK-NEXT: ret;
%r = call float @llvm.convert.from.arbitrary.fp.f32.i8(i8 124, metadata !"Float8E5M2")
ret float %r
}
; Float8E5M2 NaN: 0_11111_01
define float @from_f8e5m2_nan() {
; CHECK-LABEL: from_f8e5m2_nan(
; CHECK: {
; CHECK-EMPTY:
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: st.param.b32 [func_retval0], 2143289344;
; CHECK-NEXT: ret;
%r = call float @llvm.convert.from.arbitrary.fp.f32.i8(i8 125, metadata !"Float8E5M2")
ret float %r
}
; Float8E5M2 max: 0_11110_11 = 57344
define float @from_f8e5m2_max() {
; CHECK-LABEL: from_f8e5m2_max(
; CHECK: {
; CHECK-EMPTY:
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: st.param.b32 [func_retval0], 1197473792;
; CHECK-NEXT: ret;
%r = call float @llvm.convert.from.arbitrary.fp.f32.i8(i8 123, metadata !"Float8E5M2")
ret float %r
}
; Float8E5M2 negative: 1_01111_00 = -1.0
define float @from_f8e5m2_neg() {
; CHECK-LABEL: from_f8e5m2_neg(
; CHECK: {
; CHECK-EMPTY:
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: st.param.b32 [func_retval0], -1082130432;
; CHECK-NEXT: ret;
%r = call float @llvm.convert.from.arbitrary.fp.f32.i8(i8 -68, metadata !"Float8E5M2")
ret float %r
}
; Float8E5M2 runtime arg test
define float @from_f8e5m2_dynamic(i8 %x) {
; CHECK-LABEL: from_f8e5m2_dynamic(
; CHECK: {
; CHECK-NEXT: .reg .pred %p<6>;
; CHECK-NEXT: .reg .b32 %r<31>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b8 %r1, [from_f8e5m2_dynamic_param_0];
; CHECK-NEXT: shl.b32 %r2, %r1, 24;
; CHECK-NEXT: and.b32 %r3, %r2, -2147483648;
; CHECK-NEXT: and.b32 %r4, %r1, 3;
; CHECK-NEXT: clz.b32 %r5, %r4;
; CHECK-NEXT: sub.s32 %r6, 142, %r5;
; CHECK-NEXT: shl.b32 %r7, %r6, 23;
; CHECK-NEXT: or.b32 %r8, %r3, %r7;
; CHECK-NEXT: sub.s32 %r9, 31, %r5;
; CHECK-NEXT: mov.b32 %r10, 1;
; CHECK-NEXT: shl.b32 %r11, %r10, %r9;
; CHECK-NEXT: xor.b32 %r12, %r4, %r11;
; CHECK-NEXT: add.s32 %r13, %r5, -8;
; CHECK-NEXT: shl.b32 %r14, %r12, %r13;
; CHECK-NEXT: or.b32 %r15, %r8, %r14;
; CHECK-NEXT: bfe.u32 %r16, %r1, 2, 5;
; CHECK-NEXT: shl.b32 %r17, %r16, 23;
; CHECK-NEXT: or.b32 %r18, %r17, %r3;
; CHECK-NEXT: shl.b32 %r19, %r4, 21;
; CHECK-NEXT: or.b32 %r20, %r18, %r19;
; CHECK-NEXT: add.s32 %r21, %r20, 939524096;
; CHECK-NEXT: setp.ne.b32 %p1, %r4, 0;
; CHECK-NEXT: selp.b32 %r22, %r15, %r21, %p1;
; CHECK-NEXT: setp.eq.b32 %p2, %r16, 0;
; CHECK-NEXT: selp.b32 %r23, %r22, %r21, %p2;
; CHECK-NEXT: or.b32 %r24, %r16, %r4;
; CHECK-NEXT: setp.eq.b32 %p3, %r24, 0;
; CHECK-NEXT: selp.b32 %r25, %r3, %r23, %p3;
; CHECK-NEXT: setp.eq.b32 %p4, %r4, 0;
; CHECK-NEXT: or.b32 %r26, %r3, 2139095040;
; CHECK-NEXT: selp.b32 %r27, %r26, %r25, %p4;
; CHECK-NEXT: setp.eq.b32 %p5, %r16, 31;
; CHECK-NEXT: selp.b32 %r28, %r27, %r25, %p5;
; CHECK-NEXT: selp.b32 %r29, 2143289344, %r28, %p1;
; CHECK-NEXT: selp.b32 %r30, %r29, %r28, %p5;
; CHECK-NEXT: st.param.b32 [func_retval0], %r30;
; CHECK-NEXT: ret;
%r = call float @llvm.convert.from.arbitrary.fp.f32.i8(i8 %x, metadata !"Float8E5M2")
ret float %r
}
; Float8E4M3FN (NanOnly, NanEncoding=AllOnes)
; Layout: sign(1) exp(4) mant(3), maxExp=8, minExp=-6, bias=7
; Only 0_1111_111 and 1_1111_111 are NaN; all other exp=15 values are finite.
; Float8E4M3FN normal: 0_0111_000 = 1.0
define float @from_f8e4m3fn_normal() {
; CHECK-LABEL: from_f8e4m3fn_normal(
; CHECK: {
; CHECK-EMPTY:
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: st.param.b32 [func_retval0], 1065353216;
; CHECK-NEXT: ret;
%r = call float @llvm.convert.from.arbitrary.fp.f32.i8(i8 56, metadata !"Float8E4M3FN")
ret float %r
}
; Float8E4M3FN NaN: 0_1111_111
define float @from_f8e4m3fn_nan() {
; CHECK-LABEL: from_f8e4m3fn_nan(
; CHECK: {
; CHECK-EMPTY:
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: st.param.b32 [func_retval0], 2143289344;
; CHECK-NEXT: ret;
%r = call float @llvm.convert.from.arbitrary.fp.f32.i8(i8 127, metadata !"Float8E4M3FN")
ret float %r
}
; Float8E4M3FN not-NaN: 0_1111_110 = 448
; Despite exp=all-ones, this is a valid finite number (max value)
define float @from_f8e4m3fn_max() {
; CHECK-LABEL: from_f8e4m3fn_max(
; CHECK: {
; CHECK-EMPTY:
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: st.param.b32 [func_retval0], 1138753536;
; CHECK-NEXT: ret;
%r = call float @llvm.convert.from.arbitrary.fp.f32.i8(i8 126, metadata !"Float8E4M3FN")
ret float %r
}
; Float8E4M3FN not-NaN: 0_1111_101 = 416
; exp=all-ones but mant!=all-ones so this is finite
define float @from_f8e4m3fn_not_nan() {
; CHECK-LABEL: from_f8e4m3fn_not_nan(
; CHECK: {
; CHECK-EMPTY:
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: st.param.b32 [func_retval0], 1137704960;
; CHECK-NEXT: ret;
%r = call float @llvm.convert.from.arbitrary.fp.f32.i8(i8 125, metadata !"Float8E4M3FN")
ret float %r
}
; Float8E4M3FN zero: 0_0000_000 = +0.0
define float @from_f8e4m3fn_zero() {
; CHECK-LABEL: from_f8e4m3fn_zero(
; CHECK: {
; CHECK-EMPTY:
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: st.param.b32 [func_retval0], 0;
; CHECK-NEXT: ret;
%r = call float @llvm.convert.from.arbitrary.fp.f32.i8(i8 0, metadata !"Float8E4M3FN")
ret float %r
}
; Float8E4M3FN denorm: 0_0000_001 = 2^(-9)
define float @from_f8e4m3fn_denorm() {
; CHECK-LABEL: from_f8e4m3fn_denorm(
; CHECK: {
; CHECK-EMPTY:
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: st.param.b32 [func_retval0], 989855744;
; CHECK-NEXT: ret;
%r = call float @llvm.convert.from.arbitrary.fp.f32.i8(i8 1, metadata !"Float8E4M3FN")
ret float %r
}
; Float8E4M3FN runtime arg test
define float @from_f8e4m3fn_dynamic(i8 %x) {
; CHECK-LABEL: from_f8e4m3fn_dynamic(
; CHECK: {
; CHECK-NEXT: .reg .pred %p<6>;
; CHECK-NEXT: .reg .b32 %r<28>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b8 %r1, [from_f8e4m3fn_dynamic_param_0];
; CHECK-NEXT: shl.b32 %r2, %r1, 24;
; CHECK-NEXT: and.b32 %r3, %r2, -2147483648;
; CHECK-NEXT: and.b32 %r4, %r1, 7;
; CHECK-NEXT: clz.b32 %r5, %r4;
; CHECK-NEXT: sub.s32 %r6, 149, %r5;
; CHECK-NEXT: shl.b32 %r7, %r6, 23;
; CHECK-NEXT: or.b32 %r8, %r3, %r7;
; CHECK-NEXT: sub.s32 %r9, 31, %r5;
; CHECK-NEXT: mov.b32 %r10, 1;
; CHECK-NEXT: shl.b32 %r11, %r10, %r9;
; CHECK-NEXT: xor.b32 %r12, %r4, %r11;
; CHECK-NEXT: add.s32 %r13, %r5, -8;
; CHECK-NEXT: shl.b32 %r14, %r12, %r13;
; CHECK-NEXT: or.b32 %r15, %r8, %r14;
; CHECK-NEXT: bfe.u32 %r16, %r1, 3, 4;
; CHECK-NEXT: shl.b32 %r17, %r16, 23;
; CHECK-NEXT: or.b32 %r18, %r17, %r3;
; CHECK-NEXT: shl.b32 %r19, %r4, 20;
; CHECK-NEXT: or.b32 %r20, %r18, %r19;
; CHECK-NEXT: add.s32 %r21, %r20, 1006632960;
; CHECK-NEXT: setp.ne.b32 %p1, %r4, 0;
; CHECK-NEXT: selp.b32 %r22, %r15, %r21, %p1;
; CHECK-NEXT: setp.eq.b32 %p2, %r16, 0;
; CHECK-NEXT: selp.b32 %r23, %r22, %r21, %p2;
; CHECK-NEXT: or.b32 %r24, %r16, %r4;
; CHECK-NEXT: setp.eq.b32 %p3, %r24, 0;
; CHECK-NEXT: selp.b32 %r25, %r3, %r23, %p3;
; CHECK-NEXT: setp.eq.b32 %p4, %r4, 7;
; CHECK-NEXT: selp.b32 %r26, 2143289344, %r25, %p4;
; CHECK-NEXT: setp.eq.b32 %p5, %r16, 15;
; CHECK-NEXT: selp.b32 %r27, %r26, %r25, %p5;
; CHECK-NEXT: st.param.b32 [func_retval0], %r27;
; CHECK-NEXT: ret;
%r = call float @llvm.convert.from.arbitrary.fp.f32.i8(i8 %x, metadata !"Float8E4M3FN")
ret float %r
}
; Float6E3M2FN (FiniteOnly)
; Layout: sign(1) exp(3) mant(2), bias=3, maxExp=4
; No Inf, no NaN. All bit patterns are finite.
; Float6E3M2FN normal: 0_011_00 = 1.0
define float @from_f6e3m2fn_normal() {
; CHECK-LABEL: from_f6e3m2fn_normal(
; CHECK: {
; CHECK-EMPTY:
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: st.param.b32 [func_retval0], 1065353216;
; CHECK-NEXT: ret;
%r = call float @llvm.convert.from.arbitrary.fp.f32.i6(i6 12, metadata !"Float6E3M2FN")
ret float %r
}
; Float6E3M2FN max: 0_111_11 = 28.0
define float @from_f6e3m2fn_max() {
; CHECK-LABEL: from_f6e3m2fn_max(
; CHECK: {
; CHECK-EMPTY:
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: st.param.b32 [func_retval0], 1105199104;
; CHECK-NEXT: ret;
%r = call float @llvm.convert.from.arbitrary.fp.f32.i6(i6 31, metadata !"Float6E3M2FN")
ret float %r
}
; Float6E3M2FN denorm: 0_000_01 = 0.0625
define float @from_f6e3m2fn_denorm() {
; CHECK-LABEL: from_f6e3m2fn_denorm(
; CHECK: {
; CHECK-EMPTY:
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: st.param.b32 [func_retval0], 1031798784;
; CHECK-NEXT: ret;
%r = call float @llvm.convert.from.arbitrary.fp.f32.i6(i6 1, metadata !"Float6E3M2FN")
ret float %r
}
; Float6E3M2FN zero: 0_000_00 = +0.0
define float @from_f6e3m2fn_zero() {
; CHECK-LABEL: from_f6e3m2fn_zero(
; CHECK: {
; CHECK-EMPTY:
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: st.param.b32 [func_retval0], 0;
; CHECK-NEXT: ret;
%r = call float @llvm.convert.from.arbitrary.fp.f32.i6(i6 0, metadata !"Float6E3M2FN")
ret float %r
}
; Float6E3M2FN negative: 1_011_00 = -1.0
define float @from_f6e3m2fn_neg() {
; CHECK-LABEL: from_f6e3m2fn_neg(
; CHECK: {
; CHECK-EMPTY:
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: st.param.b32 [func_retval0], -1082130432;
; CHECK-NEXT: ret;
%r = call float @llvm.convert.from.arbitrary.fp.f32.i6(i6 -20, metadata !"Float6E3M2FN")
ret float %r
}
; Float6E3M2FN runtime arg test
define float @from_f6e3m2fn_dynamic(i6 %x) {
; CHECK-LABEL: from_f6e3m2fn_dynamic(
; CHECK: {
; CHECK-NEXT: .reg .pred %p<4>;
; CHECK-NEXT: .reg .b16 %rs<5>;
; CHECK-NEXT: .reg .b32 %r<26>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b8 %rs1, [from_f6e3m2fn_dynamic_param_0+1];
; CHECK-NEXT: shl.b16 %rs2, %rs1, 8;
; CHECK-NEXT: ld.param.b8 %rs3, [from_f6e3m2fn_dynamic_param_0];
; CHECK-NEXT: or.b16 %rs4, %rs2, %rs3;
; CHECK-NEXT: cvt.u32.u16 %r1, %rs4;
; CHECK-NEXT: shl.b32 %r2, %r1, 26;
; CHECK-NEXT: and.b32 %r3, %r2, -2147483648;
; CHECK-NEXT: and.b32 %r4, %r1, 3;
; CHECK-NEXT: clz.b32 %r5, %r4;
; CHECK-NEXT: sub.s32 %r6, 154, %r5;
; CHECK-NEXT: shl.b32 %r7, %r6, 23;
; CHECK-NEXT: or.b32 %r8, %r3, %r7;
; CHECK-NEXT: sub.s32 %r9, 31, %r5;
; CHECK-NEXT: mov.b32 %r10, 1;
; CHECK-NEXT: shl.b32 %r11, %r10, %r9;
; CHECK-NEXT: xor.b32 %r12, %r4, %r11;
; CHECK-NEXT: add.s32 %r13, %r5, -8;
; CHECK-NEXT: shl.b32 %r14, %r12, %r13;
; CHECK-NEXT: or.b32 %r15, %r8, %r14;
; CHECK-NEXT: bfe.u32 %r16, %r1, 2, 3;
; CHECK-NEXT: shl.b32 %r17, %r16, 23;
; CHECK-NEXT: or.b32 %r18, %r17, %r3;
; CHECK-NEXT: shl.b32 %r19, %r4, 21;
; CHECK-NEXT: or.b32 %r20, %r18, %r19;
; CHECK-NEXT: add.s32 %r21, %r20, 1040187392;
; CHECK-NEXT: setp.ne.b32 %p1, %r4, 0;
; CHECK-NEXT: selp.b32 %r22, %r15, %r21, %p1;
; CHECK-NEXT: setp.eq.b32 %p2, %r16, 0;
; CHECK-NEXT: selp.b32 %r23, %r22, %r21, %p2;
; CHECK-NEXT: or.b32 %r24, %r16, %r4;
; CHECK-NEXT: setp.eq.b32 %p3, %r24, 0;
; CHECK-NEXT: selp.b32 %r25, %r3, %r23, %p3;
; CHECK-NEXT: st.param.b32 [func_retval0], %r25;
; CHECK-NEXT: ret;
%r = call float @llvm.convert.from.arbitrary.fp.f32.i6(i6 %x, metadata !"Float6E3M2FN")
ret float %r
}
; Float6E2M3FN (FiniteOnly)
; Layout: sign(1) exp(2) mant(3), bias=1, maxExp=2
; No Inf, no NaN. All bit patterns are finite.
; Float6E2M3FN normal: 0_01_000 = 1.0
define float @from_f6e2m3fn_normal() {
; CHECK-LABEL: from_f6e2m3fn_normal(
; CHECK: {
; CHECK-EMPTY:
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: st.param.b32 [func_retval0], 1065353216;
; CHECK-NEXT: ret;
%r = call float @llvm.convert.from.arbitrary.fp.f32.i6(i6 8, metadata !"Float6E2M3FN")
ret float %r
}
; Float6E2M3FN max: 0_11_111 = 7.5
define float @from_f6e2m3fn_max() {
; CHECK-LABEL: from_f6e2m3fn_max(
; CHECK: {
; CHECK-EMPTY:
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: st.param.b32 [func_retval0], 1089470464;
; CHECK-NEXT: ret;
%r = call float @llvm.convert.from.arbitrary.fp.f32.i6(i6 31, metadata !"Float6E2M3FN")
ret float %r
}
; Float6E2M3FN denorm: 0_00_001 = 0.125
define float @from_f6e2m3fn_denorm() {
; CHECK-LABEL: from_f6e2m3fn_denorm(
; CHECK: {
; CHECK-EMPTY:
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: st.param.b32 [func_retval0], 1040187392;
; CHECK-NEXT: ret;
%r = call float @llvm.convert.from.arbitrary.fp.f32.i6(i6 1, metadata !"Float6E2M3FN")
ret float %r
}
; Float6E2M3FN zero: 0_00_000 = +0.0
define float @from_f6e2m3fn_zero() {
; CHECK-LABEL: from_f6e2m3fn_zero(
; CHECK: {
; CHECK-EMPTY:
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: st.param.b32 [func_retval0], 0;
; CHECK-NEXT: ret;
%r = call float @llvm.convert.from.arbitrary.fp.f32.i6(i6 0, metadata !"Float6E2M3FN")
ret float %r
}
; Float6E2M3FN runtime arg test
define float @from_f6e2m3fn_dynamic(i6 %x) {
; CHECK-LABEL: from_f6e2m3fn_dynamic(
; CHECK: {
; CHECK-NEXT: .reg .pred %p<4>;
; CHECK-NEXT: .reg .b16 %rs<5>;
; CHECK-NEXT: .reg .b32 %r<26>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b8 %rs1, [from_f6e2m3fn_dynamic_param_0+1];
; CHECK-NEXT: shl.b16 %rs2, %rs1, 8;
; CHECK-NEXT: ld.param.b8 %rs3, [from_f6e2m3fn_dynamic_param_0];
; CHECK-NEXT: or.b16 %rs4, %rs2, %rs3;
; CHECK-NEXT: cvt.u32.u16 %r1, %rs4;
; CHECK-NEXT: shl.b32 %r2, %r1, 26;
; CHECK-NEXT: and.b32 %r3, %r2, -2147483648;
; CHECK-NEXT: and.b32 %r4, %r1, 7;
; CHECK-NEXT: clz.b32 %r5, %r4;
; CHECK-NEXT: sub.s32 %r6, 155, %r5;
; CHECK-NEXT: shl.b32 %r7, %r6, 23;
; CHECK-NEXT: or.b32 %r8, %r3, %r7;
; CHECK-NEXT: sub.s32 %r9, 31, %r5;
; CHECK-NEXT: mov.b32 %r10, 1;
; CHECK-NEXT: shl.b32 %r11, %r10, %r9;
; CHECK-NEXT: xor.b32 %r12, %r4, %r11;
; CHECK-NEXT: add.s32 %r13, %r5, -8;
; CHECK-NEXT: shl.b32 %r14, %r12, %r13;
; CHECK-NEXT: or.b32 %r15, %r8, %r14;
; CHECK-NEXT: bfe.u32 %r16, %r1, 3, 2;
; CHECK-NEXT: shl.b32 %r17, %r16, 23;
; CHECK-NEXT: or.b32 %r18, %r17, %r3;
; CHECK-NEXT: shl.b32 %r19, %r4, 20;
; CHECK-NEXT: or.b32 %r20, %r18, %r19;
; CHECK-NEXT: add.s32 %r21, %r20, 1056964608;
; CHECK-NEXT: setp.ne.b32 %p1, %r4, 0;
; CHECK-NEXT: selp.b32 %r22, %r15, %r21, %p1;
; CHECK-NEXT: setp.eq.b32 %p2, %r16, 0;
; CHECK-NEXT: selp.b32 %r23, %r22, %r21, %p2;
; CHECK-NEXT: or.b32 %r24, %r16, %r4;
; CHECK-NEXT: setp.eq.b32 %p3, %r24, 0;
; CHECK-NEXT: selp.b32 %r25, %r3, %r23, %p3;
; CHECK-NEXT: st.param.b32 [func_retval0], %r25;
; CHECK-NEXT: ret;
%r = call float @llvm.convert.from.arbitrary.fp.f32.i6(i6 %x, metadata !"Float6E2M3FN")
ret float %r
}
; Float4E2M1FN (FiniteOnly)
; Layout: sign(1) exp(2) mant(1), bias=1, maxExp=2
; No Inf, no NaN.
; Float4E2M1FN normal: 0_01_0 = 1.0
define float @from_f4e2m1fn_normal() {
; CHECK-LABEL: from_f4e2m1fn_normal(
; CHECK: {
; CHECK-EMPTY:
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: st.param.b32 [func_retval0], 1065353216;
; CHECK-NEXT: ret;
%r = call float @llvm.convert.from.arbitrary.fp.f32.i4(i4 2, metadata !"Float4E2M1FN")
ret float %r
}
; Float4E2M1FN denorm: 0_00_1 = 0.5
define float @from_f4e2m1fn_denorm() {
; CHECK-LABEL: from_f4e2m1fn_denorm(
; CHECK: {
; CHECK-EMPTY:
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: st.param.b32 [func_retval0], 1056964608;
; CHECK-NEXT: ret;
%r = call float @llvm.convert.from.arbitrary.fp.f32.i4(i4 1, metadata !"Float4E2M1FN")
ret float %r
}
; Float4E2M1FN max: 0_11_1 = 6.0
define float @from_f4e2m1fn_max() {
; CHECK-LABEL: from_f4e2m1fn_max(
; CHECK: {
; CHECK-EMPTY:
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: st.param.b32 [func_retval0], 1086324736;
; CHECK-NEXT: ret;
%r = call float @llvm.convert.from.arbitrary.fp.f32.i4(i4 7, metadata !"Float4E2M1FN")
ret float %r
}
; Float4E2M1FN runtime arg test
define float @from_f4e2m1fn_dynamic(i4 %x) {
; CHECK-LABEL: from_f4e2m1fn_dynamic(
; CHECK: {
; CHECK-NEXT: .reg .pred %p<4>;
; CHECK-NEXT: .reg .b16 %rs<6>;
; CHECK-NEXT: .reg .b32 %r<26>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b8 %rs1, [from_f4e2m1fn_dynamic_param_0+1];
; CHECK-NEXT: shl.b16 %rs2, %rs1, 8;
; CHECK-NEXT: ld.param.b8 %rs3, [from_f4e2m1fn_dynamic_param_0];
; CHECK-NEXT: or.b16 %rs4, %rs2, %rs3;
; CHECK-NEXT: cvt.u32.u16 %r1, %rs4;
; CHECK-NEXT: shl.b32 %r2, %r1, 28;
; CHECK-NEXT: and.b32 %r3, %r2, -2147483648;
; CHECK-NEXT: and.b32 %r4, %r1, 1;
; CHECK-NEXT: clz.b32 %r5, %r4;
; CHECK-NEXT: sub.s32 %r6, 157, %r5;
; CHECK-NEXT: shl.b32 %r7, %r6, 23;
; CHECK-NEXT: or.b32 %r8, %r3, %r7;
; CHECK-NEXT: sub.s32 %r9, 31, %r5;
; CHECK-NEXT: mov.b32 %r10, 1;
; CHECK-NEXT: shl.b32 %r11, %r10, %r9;
; CHECK-NEXT: xor.b32 %r12, %r4, %r11;
; CHECK-NEXT: add.s32 %r13, %r5, -8;
; CHECK-NEXT: shl.b32 %r14, %r12, %r13;
; CHECK-NEXT: or.b32 %r15, %r8, %r14;
; CHECK-NEXT: bfe.u32 %r16, %r1, 1, 2;
; CHECK-NEXT: shl.b32 %r17, %r16, 23;
; CHECK-NEXT: or.b32 %r18, %r17, %r3;
; CHECK-NEXT: shl.b32 %r19, %r4, 22;
; CHECK-NEXT: or.b32 %r20, %r18, %r19;
; CHECK-NEXT: add.s32 %r21, %r20, 1056964608;
; CHECK-NEXT: and.b16 %rs5, %rs3, 1;
; CHECK-NEXT: setp.ne.b16 %p1, %rs5, 0;
; CHECK-NEXT: selp.b32 %r22, %r15, %r21, %p1;
; CHECK-NEXT: setp.eq.b32 %p2, %r16, 0;
; CHECK-NEXT: selp.b32 %r23, %r22, %r21, %p2;
; CHECK-NEXT: or.b32 %r24, %r16, %r4;
; CHECK-NEXT: setp.eq.b32 %p3, %r24, 0;
; CHECK-NEXT: selp.b32 %r25, %r3, %r23, %p3;
; CHECK-NEXT: st.param.b32 [func_retval0], %r25;
; CHECK-NEXT: ret;
%r = call float @llvm.convert.from.arbitrary.fp.f32.i4(i4 %x, metadata !"Float4E2M1FN")
ret float %r
}
; Float8E5M2 to f16: 1.0
define half @from_f8e5m2_to_f16() {
; CHECK-LABEL: from_f8e5m2_to_f16(
; CHECK: {
; CHECK-EMPTY:
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: st.param.b16 [func_retval0], 0x3C00;
; CHECK-NEXT: ret;
%r = call half @llvm.convert.from.arbitrary.fp.f16.i8(i8 60, metadata !"Float8E5M2")
ret half %r
}
; Float8E5M2 to f64: 1.0
define double @from_f8e5m2_to_f64() {
; CHECK-LABEL: from_f8e5m2_to_f64(
; CHECK: {
; CHECK-EMPTY:
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: st.param.b64 [func_retval0], 4607182418800017408;
; CHECK-NEXT: ret;
%r = call double @llvm.convert.from.arbitrary.fp.f64.i8(i8 60, metadata !"Float8E5M2")
ret double %r
}
; Vector test: Float4E2M1FN <4 x i4> -> <4 x float>
define <4 x float> @fp4_to_f32_vec(<4 x i4> %x) {
; CHECK-LABEL: fp4_to_f32_vec(
; CHECK: {
; CHECK-NEXT: .reg .pred %p<13>;
; CHECK-NEXT: .reg .b32 %r<101>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %r1, [fp4_to_f32_vec_param_0+2];
; CHECK-NEXT: shl.b32 %r2, %r1, 16;
; CHECK-NEXT: ld.param.b16 %r3, [fp4_to_f32_vec_param_0];
; CHECK-NEXT: prmt.b32 %r4, %r3, 0, 0x7771U;
; CHECK-NEXT: shl.b32 %r5, %r4, 28;
; CHECK-NEXT: and.b32 %r6, %r5, -2147483648;
; CHECK-NEXT: and.b32 %r7, %r4, 1;
; CHECK-NEXT: clz.b32 %r8, %r7;
; CHECK-NEXT: sub.s32 %r9, 157, %r8;
; CHECK-NEXT: shl.b32 %r10, %r9, 23;
; CHECK-NEXT: or.b32 %r11, %r6, %r10;
; CHECK-NEXT: sub.s32 %r12, 31, %r8;
; CHECK-NEXT: mov.b32 %r13, 1;
; CHECK-NEXT: shl.b32 %r14, %r13, %r12;
; CHECK-NEXT: xor.b32 %r15, %r7, %r14;
; CHECK-NEXT: add.s32 %r16, %r8, -8;
; CHECK-NEXT: shl.b32 %r17, %r15, %r16;
; CHECK-NEXT: or.b32 %r18, %r11, %r17;
; CHECK-NEXT: bfe.u32 %r19, %r4, 1, 2;
; CHECK-NEXT: shl.b32 %r20, %r19, 23;
; CHECK-NEXT: or.b32 %r21, %r20, %r6;
; CHECK-NEXT: shl.b32 %r22, %r7, 22;
; CHECK-NEXT: or.b32 %r23, %r21, %r22;
; CHECK-NEXT: add.s32 %r24, %r23, 1056964608;
; CHECK-NEXT: setp.ne.b32 %p1, %r7, 0;
; CHECK-NEXT: selp.b32 %r25, %r18, %r24, %p1;
; CHECK-NEXT: setp.eq.b32 %p2, %r19, 0;
; CHECK-NEXT: selp.b32 %r26, %r25, %r24, %p2;
; CHECK-NEXT: or.b32 %r27, %r19, %r7;
; CHECK-NEXT: setp.eq.b32 %p3, %r27, 0;
; CHECK-NEXT: selp.b32 %r28, %r6, %r26, %p3;
; CHECK-NEXT: prmt.b32 %r29, %r3, 0, 0x7770U;
; CHECK-NEXT: shl.b32 %r30, %r29, 28;
; CHECK-NEXT: and.b32 %r31, %r30, -2147483648;
; CHECK-NEXT: and.b32 %r32, %r29, 1;
; CHECK-NEXT: clz.b32 %r33, %r32;
; CHECK-NEXT: sub.s32 %r34, 157, %r33;
; CHECK-NEXT: shl.b32 %r35, %r34, 23;
; CHECK-NEXT: or.b32 %r36, %r31, %r35;
; CHECK-NEXT: sub.s32 %r37, 31, %r33;
; CHECK-NEXT: shl.b32 %r38, %r13, %r37;
; CHECK-NEXT: xor.b32 %r39, %r32, %r38;
; CHECK-NEXT: add.s32 %r40, %r33, -8;
; CHECK-NEXT: shl.b32 %r41, %r39, %r40;
; CHECK-NEXT: or.b32 %r42, %r36, %r41;
; CHECK-NEXT: bfe.u32 %r43, %r29, 1, 2;
; CHECK-NEXT: shl.b32 %r44, %r43, 23;
; CHECK-NEXT: or.b32 %r45, %r44, %r31;
; CHECK-NEXT: shl.b32 %r46, %r32, 22;
; CHECK-NEXT: or.b32 %r47, %r45, %r46;
; CHECK-NEXT: add.s32 %r48, %r47, 1056964608;
; CHECK-NEXT: setp.ne.b32 %p4, %r32, 0;
; CHECK-NEXT: selp.b32 %r49, %r42, %r48, %p4;
; CHECK-NEXT: setp.eq.b32 %p5, %r43, 0;
; CHECK-NEXT: selp.b32 %r50, %r49, %r48, %p5;
; CHECK-NEXT: or.b32 %r51, %r43, %r32;
; CHECK-NEXT: setp.eq.b32 %p6, %r51, 0;
; CHECK-NEXT: selp.b32 %r52, %r31, %r50, %p6;
; CHECK-NEXT: prmt.b32 %r53, %r2, 0, 0x7773U;
; CHECK-NEXT: shl.b32 %r54, %r53, 28;
; CHECK-NEXT: and.b32 %r55, %r54, -2147483648;
; CHECK-NEXT: and.b32 %r56, %r53, 1;
; CHECK-NEXT: clz.b32 %r57, %r56;
; CHECK-NEXT: sub.s32 %r58, 157, %r57;
; CHECK-NEXT: shl.b32 %r59, %r58, 23;
; CHECK-NEXT: or.b32 %r60, %r55, %r59;
; CHECK-NEXT: sub.s32 %r61, 31, %r57;
; CHECK-NEXT: shl.b32 %r62, %r13, %r61;
; CHECK-NEXT: xor.b32 %r63, %r56, %r62;
; CHECK-NEXT: add.s32 %r64, %r57, -8;
; CHECK-NEXT: shl.b32 %r65, %r63, %r64;
; CHECK-NEXT: or.b32 %r66, %r60, %r65;
; CHECK-NEXT: bfe.u32 %r67, %r53, 1, 2;
; CHECK-NEXT: shl.b32 %r68, %r67, 23;
; CHECK-NEXT: or.b32 %r69, %r68, %r55;
; CHECK-NEXT: shl.b32 %r70, %r56, 22;
; CHECK-NEXT: or.b32 %r71, %r69, %r70;
; CHECK-NEXT: add.s32 %r72, %r71, 1056964608;
; CHECK-NEXT: setp.ne.b32 %p7, %r56, 0;
; CHECK-NEXT: selp.b32 %r73, %r66, %r72, %p7;
; CHECK-NEXT: setp.eq.b32 %p8, %r67, 0;
; CHECK-NEXT: selp.b32 %r74, %r73, %r72, %p8;
; CHECK-NEXT: or.b32 %r75, %r67, %r56;
; CHECK-NEXT: setp.eq.b32 %p9, %r75, 0;
; CHECK-NEXT: selp.b32 %r76, %r55, %r74, %p9;
; CHECK-NEXT: prmt.b32 %r77, %r2, 0, 0x7772U;
; CHECK-NEXT: shl.b32 %r78, %r77, 28;
; CHECK-NEXT: and.b32 %r79, %r78, -2147483648;
; CHECK-NEXT: and.b32 %r80, %r77, 1;
; CHECK-NEXT: clz.b32 %r81, %r80;
; CHECK-NEXT: sub.s32 %r82, 157, %r81;
; CHECK-NEXT: shl.b32 %r83, %r82, 23;
; CHECK-NEXT: or.b32 %r84, %r79, %r83;
; CHECK-NEXT: sub.s32 %r85, 31, %r81;
; CHECK-NEXT: shl.b32 %r86, %r13, %r85;
; CHECK-NEXT: xor.b32 %r87, %r80, %r86;
; CHECK-NEXT: add.s32 %r88, %r81, -8;
; CHECK-NEXT: shl.b32 %r89, %r87, %r88;
; CHECK-NEXT: or.b32 %r90, %r84, %r89;
; CHECK-NEXT: bfe.u32 %r91, %r77, 1, 2;
; CHECK-NEXT: shl.b32 %r92, %r91, 23;
; CHECK-NEXT: or.b32 %r93, %r92, %r79;
; CHECK-NEXT: shl.b32 %r94, %r80, 22;
; CHECK-NEXT: or.b32 %r95, %r93, %r94;
; CHECK-NEXT: add.s32 %r96, %r95, 1056964608;
; CHECK-NEXT: setp.ne.b32 %p10, %r80, 0;
; CHECK-NEXT: selp.b32 %r97, %r90, %r96, %p10;
; CHECK-NEXT: setp.eq.b32 %p11, %r91, 0;
; CHECK-NEXT: selp.b32 %r98, %r97, %r96, %p11;
; CHECK-NEXT: or.b32 %r99, %r91, %r80;
; CHECK-NEXT: setp.eq.b32 %p12, %r99, 0;
; CHECK-NEXT: selp.b32 %r100, %r79, %r98, %p12;
; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r52, %r28, %r100, %r76};
; CHECK-NEXT: ret;
%r = call <4 x float> @llvm.convert.from.arbitrary.fp.v4f32.v4i4(<4 x i4> %x, metadata !"Float4E2M1FN")
ret <4 x float> %r
}
; <2 x i8> -> <2 x float> Float8E4M3FN
define <2 x float> @from_f8e4m3fn_v2f32(<2 x i8> %x) {
; CHECK-LABEL: from_f8e4m3fn_v2f32(
; CHECK: {
; CHECK-NEXT: .reg .pred %p<11>;
; CHECK-NEXT: .reg .b16 %rs<3>;
; CHECK-NEXT: .reg .b32 %r<54>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b8 {%rs1, %rs2}, [from_f8e4m3fn_v2f32_param_0];
; CHECK-NEXT: cvt.u32.u16 %r1, %rs2;
; CHECK-NEXT: shl.b32 %r2, %r1, 24;
; CHECK-NEXT: and.b32 %r3, %r2, -2147483648;
; CHECK-NEXT: and.b32 %r4, %r1, 7;
; CHECK-NEXT: clz.b32 %r5, %r4;
; CHECK-NEXT: sub.s32 %r6, 149, %r5;
; CHECK-NEXT: shl.b32 %r7, %r6, 23;
; CHECK-NEXT: or.b32 %r8, %r3, %r7;
; CHECK-NEXT: sub.s32 %r9, 31, %r5;
; CHECK-NEXT: mov.b32 %r10, 1;
; CHECK-NEXT: shl.b32 %r11, %r10, %r9;
; CHECK-NEXT: xor.b32 %r12, %r4, %r11;
; CHECK-NEXT: add.s32 %r13, %r5, -8;
; CHECK-NEXT: shl.b32 %r14, %r12, %r13;
; CHECK-NEXT: or.b32 %r15, %r8, %r14;
; CHECK-NEXT: bfe.u32 %r16, %r1, 3, 4;
; CHECK-NEXT: shl.b32 %r17, %r16, 23;
; CHECK-NEXT: or.b32 %r18, %r17, %r3;
; CHECK-NEXT: shl.b32 %r19, %r4, 20;
; CHECK-NEXT: or.b32 %r20, %r18, %r19;
; CHECK-NEXT: add.s32 %r21, %r20, 1006632960;
; CHECK-NEXT: setp.ne.b32 %p1, %r4, 0;
; CHECK-NEXT: selp.b32 %r22, %r15, %r21, %p1;
; CHECK-NEXT: setp.eq.b32 %p2, %r16, 0;
; CHECK-NEXT: selp.b32 %r23, %r22, %r21, %p2;
; CHECK-NEXT: or.b32 %r24, %r16, %r4;
; CHECK-NEXT: setp.eq.b32 %p3, %r24, 0;
; CHECK-NEXT: selp.b32 %r25, %r3, %r23, %p3;
; CHECK-NEXT: setp.eq.b32 %p4, %r4, 7;
; CHECK-NEXT: selp.b32 %r26, 2143289344, %r25, %p4;
; CHECK-NEXT: setp.eq.b32 %p5, %r16, 15;
; CHECK-NEXT: selp.b32 %r27, %r26, %r25, %p5;
; CHECK-NEXT: cvt.u32.u16 %r28, %rs1;
; CHECK-NEXT: shl.b32 %r29, %r28, 24;
; CHECK-NEXT: and.b32 %r30, %r29, -2147483648;
; CHECK-NEXT: and.b32 %r31, %r28, 7;
; CHECK-NEXT: clz.b32 %r32, %r31;
; CHECK-NEXT: sub.s32 %r33, 149, %r32;
; CHECK-NEXT: shl.b32 %r34, %r33, 23;
; CHECK-NEXT: or.b32 %r35, %r30, %r34;
; CHECK-NEXT: sub.s32 %r36, 31, %r32;
; CHECK-NEXT: shl.b32 %r37, %r10, %r36;
; CHECK-NEXT: xor.b32 %r38, %r31, %r37;
; CHECK-NEXT: add.s32 %r39, %r32, -8;
; CHECK-NEXT: shl.b32 %r40, %r38, %r39;
; CHECK-NEXT: or.b32 %r41, %r35, %r40;
; CHECK-NEXT: bfe.u32 %r42, %r28, 3, 4;
; CHECK-NEXT: shl.b32 %r43, %r42, 23;
; CHECK-NEXT: or.b32 %r44, %r43, %r30;
; CHECK-NEXT: shl.b32 %r45, %r31, 20;
; CHECK-NEXT: or.b32 %r46, %r44, %r45;
; CHECK-NEXT: add.s32 %r47, %r46, 1006632960;
; CHECK-NEXT: setp.ne.b32 %p6, %r31, 0;
; CHECK-NEXT: selp.b32 %r48, %r41, %r47, %p6;
; CHECK-NEXT: setp.eq.b32 %p7, %r42, 0;
; CHECK-NEXT: selp.b32 %r49, %r48, %r47, %p7;
; CHECK-NEXT: or.b32 %r50, %r42, %r31;
; CHECK-NEXT: setp.eq.b32 %p8, %r50, 0;
; CHECK-NEXT: selp.b32 %r51, %r30, %r49, %p8;
; CHECK-NEXT: setp.eq.b32 %p9, %r31, 7;
; CHECK-NEXT: selp.b32 %r52, 2143289344, %r51, %p9;
; CHECK-NEXT: setp.eq.b32 %p10, %r42, 15;
; CHECK-NEXT: selp.b32 %r53, %r52, %r51, %p10;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r53, %r27};
; CHECK-NEXT: ret;
%r = call <2 x float> @llvm.convert.from.arbitrary.fp.v2f32.v2i8(<2 x i8> %x, metadata !"Float8E4M3FN")
ret <2 x float> %r
}
; <2 x i8> -> <2 x float> Float8E5M2
define <2 x float> @from_f8e5m2_v2f32(<2 x i8> %x) {
; CHECK-LABEL: from_f8e5m2_v2f32(
; CHECK: {
; CHECK-NEXT: .reg .pred %p<11>;
; CHECK-NEXT: .reg .b16 %rs<3>;
; CHECK-NEXT: .reg .b32 %r<60>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b8 {%rs1, %rs2}, [from_f8e5m2_v2f32_param_0];
; CHECK-NEXT: cvt.u32.u16 %r1, %rs2;
; CHECK-NEXT: shl.b32 %r2, %r1, 24;
; CHECK-NEXT: and.b32 %r3, %r2, -2147483648;
; CHECK-NEXT: and.b32 %r4, %r1, 3;
; CHECK-NEXT: clz.b32 %r5, %r4;
; CHECK-NEXT: sub.s32 %r6, 142, %r5;
; CHECK-NEXT: shl.b32 %r7, %r6, 23;
; CHECK-NEXT: or.b32 %r8, %r3, %r7;
; CHECK-NEXT: sub.s32 %r9, 31, %r5;
; CHECK-NEXT: mov.b32 %r10, 1;
; CHECK-NEXT: shl.b32 %r11, %r10, %r9;
; CHECK-NEXT: xor.b32 %r12, %r4, %r11;
; CHECK-NEXT: add.s32 %r13, %r5, -8;
; CHECK-NEXT: shl.b32 %r14, %r12, %r13;
; CHECK-NEXT: or.b32 %r15, %r8, %r14;
; CHECK-NEXT: bfe.u32 %r16, %r1, 2, 5;
; CHECK-NEXT: shl.b32 %r17, %r16, 23;
; CHECK-NEXT: or.b32 %r18, %r17, %r3;
; CHECK-NEXT: shl.b32 %r19, %r4, 21;
; CHECK-NEXT: or.b32 %r20, %r18, %r19;
; CHECK-NEXT: add.s32 %r21, %r20, 939524096;
; CHECK-NEXT: setp.ne.b32 %p1, %r4, 0;
; CHECK-NEXT: selp.b32 %r22, %r15, %r21, %p1;
; CHECK-NEXT: setp.eq.b32 %p2, %r16, 0;
; CHECK-NEXT: selp.b32 %r23, %r22, %r21, %p2;
; CHECK-NEXT: or.b32 %r24, %r16, %r4;
; CHECK-NEXT: setp.eq.b32 %p3, %r24, 0;
; CHECK-NEXT: selp.b32 %r25, %r3, %r23, %p3;
; CHECK-NEXT: setp.eq.b32 %p4, %r4, 0;
; CHECK-NEXT: or.b32 %r26, %r3, 2139095040;
; CHECK-NEXT: selp.b32 %r27, %r26, %r25, %p4;
; CHECK-NEXT: setp.eq.b32 %p5, %r16, 31;
; CHECK-NEXT: selp.b32 %r28, %r27, %r25, %p5;
; CHECK-NEXT: selp.b32 %r29, 2143289344, %r28, %p1;
; CHECK-NEXT: selp.b32 %r30, %r29, %r28, %p5;
; CHECK-NEXT: cvt.u32.u16 %r31, %rs1;
; CHECK-NEXT: shl.b32 %r32, %r31, 24;
; CHECK-NEXT: and.b32 %r33, %r32, -2147483648;
; CHECK-NEXT: and.b32 %r34, %r31, 3;
; CHECK-NEXT: clz.b32 %r35, %r34;
; CHECK-NEXT: sub.s32 %r36, 142, %r35;
; CHECK-NEXT: shl.b32 %r37, %r36, 23;
; CHECK-NEXT: or.b32 %r38, %r33, %r37;
; CHECK-NEXT: sub.s32 %r39, 31, %r35;
; CHECK-NEXT: shl.b32 %r40, %r10, %r39;
; CHECK-NEXT: xor.b32 %r41, %r34, %r40;
; CHECK-NEXT: add.s32 %r42, %r35, -8;
; CHECK-NEXT: shl.b32 %r43, %r41, %r42;
; CHECK-NEXT: or.b32 %r44, %r38, %r43;
; CHECK-NEXT: bfe.u32 %r45, %r31, 2, 5;
; CHECK-NEXT: shl.b32 %r46, %r45, 23;
; CHECK-NEXT: or.b32 %r47, %r46, %r33;
; CHECK-NEXT: shl.b32 %r48, %r34, 21;
; CHECK-NEXT: or.b32 %r49, %r47, %r48;
; CHECK-NEXT: add.s32 %r50, %r49, 939524096;
; CHECK-NEXT: setp.ne.b32 %p6, %r34, 0;
; CHECK-NEXT: selp.b32 %r51, %r44, %r50, %p6;
; CHECK-NEXT: setp.eq.b32 %p7, %r45, 0;
; CHECK-NEXT: selp.b32 %r52, %r51, %r50, %p7;
; CHECK-NEXT: or.b32 %r53, %r45, %r34;
; CHECK-NEXT: setp.eq.b32 %p8, %r53, 0;
; CHECK-NEXT: selp.b32 %r54, %r33, %r52, %p8;
; CHECK-NEXT: setp.eq.b32 %p9, %r34, 0;
; CHECK-NEXT: or.b32 %r55, %r33, 2139095040;
; CHECK-NEXT: selp.b32 %r56, %r55, %r54, %p9;
; CHECK-NEXT: setp.eq.b32 %p10, %r45, 31;
; CHECK-NEXT: selp.b32 %r57, %r56, %r54, %p10;
; CHECK-NEXT: selp.b32 %r58, 2143289344, %r57, %p6;
; CHECK-NEXT: selp.b32 %r59, %r58, %r57, %p10;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r59, %r30};
; CHECK-NEXT: ret;
%r = call <2 x float> @llvm.convert.from.arbitrary.fp.v2f32.v2i8(<2 x i8> %x, metadata !"Float8E5M2")
ret <2 x float> %r
}
; <3 x i8> -> <3 x float> Float8E4M3FN
define <3 x float> @from_f8e4m3fn_v3f32(<3 x i8> %x) {
; CHECK-LABEL: from_f8e4m3fn_v3f32(
; CHECK: {
; CHECK-NEXT: .reg .pred %p<16>;
; CHECK-NEXT: .reg .b32 %r<81>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [from_f8e4m3fn_v3f32_param_0];
; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x7772U;
; CHECK-NEXT: shl.b32 %r3, %r2, 24;
; CHECK-NEXT: and.b32 %r4, %r3, -2147483648;
; CHECK-NEXT: and.b32 %r5, %r2, 7;
; CHECK-NEXT: clz.b32 %r6, %r5;
; CHECK-NEXT: sub.s32 %r7, 149, %r6;
; CHECK-NEXT: shl.b32 %r8, %r7, 23;
; CHECK-NEXT: or.b32 %r9, %r4, %r8;
; CHECK-NEXT: sub.s32 %r10, 31, %r6;
; CHECK-NEXT: mov.b32 %r11, 1;
; CHECK-NEXT: shl.b32 %r12, %r11, %r10;
; CHECK-NEXT: xor.b32 %r13, %r5, %r12;
; CHECK-NEXT: add.s32 %r14, %r6, -8;
; CHECK-NEXT: shl.b32 %r15, %r13, %r14;
; CHECK-NEXT: or.b32 %r16, %r9, %r15;
; CHECK-NEXT: bfe.u32 %r17, %r2, 3, 4;
; CHECK-NEXT: shl.b32 %r18, %r17, 23;
; CHECK-NEXT: or.b32 %r19, %r18, %r4;
; CHECK-NEXT: shl.b32 %r20, %r5, 20;
; CHECK-NEXT: or.b32 %r21, %r19, %r20;
; CHECK-NEXT: add.s32 %r22, %r21, 1006632960;
; CHECK-NEXT: setp.ne.b32 %p1, %r5, 0;
; CHECK-NEXT: selp.b32 %r23, %r16, %r22, %p1;
; CHECK-NEXT: setp.eq.b32 %p2, %r17, 0;
; CHECK-NEXT: selp.b32 %r24, %r23, %r22, %p2;
; CHECK-NEXT: or.b32 %r25, %r17, %r5;
; CHECK-NEXT: setp.eq.b32 %p3, %r25, 0;
; CHECK-NEXT: selp.b32 %r26, %r4, %r24, %p3;
; CHECK-NEXT: setp.eq.b32 %p4, %r5, 7;
; CHECK-NEXT: selp.b32 %r27, 2143289344, %r26, %p4;
; CHECK-NEXT: setp.eq.b32 %p5, %r17, 15;
; CHECK-NEXT: selp.b32 %r28, %r27, %r26, %p5;
; CHECK-NEXT: prmt.b32 %r29, %r1, 0, 0x7771U;
; CHECK-NEXT: shl.b32 %r30, %r29, 24;
; CHECK-NEXT: and.b32 %r31, %r30, -2147483648;
; CHECK-NEXT: and.b32 %r32, %r29, 7;
; CHECK-NEXT: clz.b32 %r33, %r32;
; CHECK-NEXT: sub.s32 %r34, 149, %r33;
; CHECK-NEXT: shl.b32 %r35, %r34, 23;
; CHECK-NEXT: or.b32 %r36, %r31, %r35;
; CHECK-NEXT: sub.s32 %r37, 31, %r33;
; CHECK-NEXT: shl.b32 %r38, %r11, %r37;
; CHECK-NEXT: xor.b32 %r39, %r32, %r38;
; CHECK-NEXT: add.s32 %r40, %r33, -8;
; CHECK-NEXT: shl.b32 %r41, %r39, %r40;
; CHECK-NEXT: or.b32 %r42, %r36, %r41;
; CHECK-NEXT: bfe.u32 %r43, %r29, 3, 4;
; CHECK-NEXT: shl.b32 %r44, %r43, 23;
; CHECK-NEXT: or.b32 %r45, %r44, %r31;
; CHECK-NEXT: shl.b32 %r46, %r32, 20;
; CHECK-NEXT: or.b32 %r47, %r45, %r46;
; CHECK-NEXT: add.s32 %r48, %r47, 1006632960;
; CHECK-NEXT: setp.ne.b32 %p6, %r32, 0;
; CHECK-NEXT: selp.b32 %r49, %r42, %r48, %p6;
; CHECK-NEXT: setp.eq.b32 %p7, %r43, 0;
; CHECK-NEXT: selp.b32 %r50, %r49, %r48, %p7;
; CHECK-NEXT: or.b32 %r51, %r43, %r32;
; CHECK-NEXT: setp.eq.b32 %p8, %r51, 0;
; CHECK-NEXT: selp.b32 %r52, %r31, %r50, %p8;
; CHECK-NEXT: setp.eq.b32 %p9, %r32, 7;
; CHECK-NEXT: selp.b32 %r53, 2143289344, %r52, %p9;
; CHECK-NEXT: setp.eq.b32 %p10, %r43, 15;
; CHECK-NEXT: selp.b32 %r54, %r53, %r52, %p10;
; CHECK-NEXT: prmt.b32 %r55, %r1, 0, 0x7770U;
; CHECK-NEXT: shl.b32 %r56, %r55, 24;
; CHECK-NEXT: and.b32 %r57, %r56, -2147483648;
; CHECK-NEXT: and.b32 %r58, %r55, 7;
; CHECK-NEXT: clz.b32 %r59, %r58;
; CHECK-NEXT: sub.s32 %r60, 149, %r59;
; CHECK-NEXT: shl.b32 %r61, %r60, 23;
; CHECK-NEXT: or.b32 %r62, %r57, %r61;
; CHECK-NEXT: sub.s32 %r63, 31, %r59;
; CHECK-NEXT: shl.b32 %r64, %r11, %r63;
; CHECK-NEXT: xor.b32 %r65, %r58, %r64;
; CHECK-NEXT: add.s32 %r66, %r59, -8;
; CHECK-NEXT: shl.b32 %r67, %r65, %r66;
; CHECK-NEXT: or.b32 %r68, %r62, %r67;
; CHECK-NEXT: bfe.u32 %r69, %r55, 3, 4;
; CHECK-NEXT: shl.b32 %r70, %r69, 23;
; CHECK-NEXT: or.b32 %r71, %r70, %r57;
; CHECK-NEXT: shl.b32 %r72, %r58, 20;
; CHECK-NEXT: or.b32 %r73, %r71, %r72;
; CHECK-NEXT: add.s32 %r74, %r73, 1006632960;
; CHECK-NEXT: setp.ne.b32 %p11, %r58, 0;
; CHECK-NEXT: selp.b32 %r75, %r68, %r74, %p11;
; CHECK-NEXT: setp.eq.b32 %p12, %r69, 0;
; CHECK-NEXT: selp.b32 %r76, %r75, %r74, %p12;
; CHECK-NEXT: or.b32 %r77, %r69, %r58;
; CHECK-NEXT: setp.eq.b32 %p13, %r77, 0;
; CHECK-NEXT: selp.b32 %r78, %r57, %r76, %p13;
; CHECK-NEXT: setp.eq.b32 %p14, %r58, 7;
; CHECK-NEXT: selp.b32 %r79, 2143289344, %r78, %p14;
; CHECK-NEXT: setp.eq.b32 %p15, %r69, 15;
; CHECK-NEXT: selp.b32 %r80, %r79, %r78, %p15;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r80, %r54};
; CHECK-NEXT: st.param.b32 [func_retval0+8], %r28;
; CHECK-NEXT: ret;
%r = call <3 x float> @llvm.convert.from.arbitrary.fp.v3f32.v3i8(<3 x i8> %x, metadata !"Float8E4M3FN")
ret <3 x float> %r
}
; <4 x i8> -> <4 x float> Float8E4M3FN
define <4 x float> @from_f8e4m3fn_v4f32(<4 x i8> %x) {
; CHECK-LABEL: from_f8e4m3fn_v4f32(
; CHECK: {
; CHECK-NEXT: .reg .pred %p<21>;
; CHECK-NEXT: .reg .b32 %r<107>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [from_f8e4m3fn_v4f32_param_0];
; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x7773U;
; CHECK-NEXT: shl.b32 %r3, %r2, 24;
; CHECK-NEXT: and.b32 %r4, %r3, -2147483648;
; CHECK-NEXT: and.b32 %r5, %r2, 7;
; CHECK-NEXT: clz.b32 %r6, %r5;
; CHECK-NEXT: sub.s32 %r7, 149, %r6;
; CHECK-NEXT: shl.b32 %r8, %r7, 23;
; CHECK-NEXT: or.b32 %r9, %r4, %r8;
; CHECK-NEXT: sub.s32 %r10, 31, %r6;
; CHECK-NEXT: mov.b32 %r11, 1;
; CHECK-NEXT: shl.b32 %r12, %r11, %r10;
; CHECK-NEXT: xor.b32 %r13, %r5, %r12;
; CHECK-NEXT: add.s32 %r14, %r6, -8;
; CHECK-NEXT: shl.b32 %r15, %r13, %r14;
; CHECK-NEXT: or.b32 %r16, %r9, %r15;
; CHECK-NEXT: bfe.u32 %r17, %r2, 3, 4;
; CHECK-NEXT: shl.b32 %r18, %r17, 23;
; CHECK-NEXT: or.b32 %r19, %r18, %r4;
; CHECK-NEXT: shl.b32 %r20, %r5, 20;
; CHECK-NEXT: or.b32 %r21, %r19, %r20;
; CHECK-NEXT: add.s32 %r22, %r21, 1006632960;
; CHECK-NEXT: setp.ne.b32 %p1, %r5, 0;
; CHECK-NEXT: selp.b32 %r23, %r16, %r22, %p1;
; CHECK-NEXT: setp.eq.b32 %p2, %r17, 0;
; CHECK-NEXT: selp.b32 %r24, %r23, %r22, %p2;
; CHECK-NEXT: or.b32 %r25, %r17, %r5;
; CHECK-NEXT: setp.eq.b32 %p3, %r25, 0;
; CHECK-NEXT: selp.b32 %r26, %r4, %r24, %p3;
; CHECK-NEXT: setp.eq.b32 %p4, %r5, 7;
; CHECK-NEXT: selp.b32 %r27, 2143289344, %r26, %p4;
; CHECK-NEXT: setp.eq.b32 %p5, %r17, 15;
; CHECK-NEXT: selp.b32 %r28, %r27, %r26, %p5;
; CHECK-NEXT: prmt.b32 %r29, %r1, 0, 0x7772U;
; CHECK-NEXT: shl.b32 %r30, %r29, 24;
; CHECK-NEXT: and.b32 %r31, %r30, -2147483648;
; CHECK-NEXT: and.b32 %r32, %r29, 7;
; CHECK-NEXT: clz.b32 %r33, %r32;
; CHECK-NEXT: sub.s32 %r34, 149, %r33;
; CHECK-NEXT: shl.b32 %r35, %r34, 23;
; CHECK-NEXT: or.b32 %r36, %r31, %r35;
; CHECK-NEXT: sub.s32 %r37, 31, %r33;
; CHECK-NEXT: shl.b32 %r38, %r11, %r37;
; CHECK-NEXT: xor.b32 %r39, %r32, %r38;
; CHECK-NEXT: add.s32 %r40, %r33, -8;
; CHECK-NEXT: shl.b32 %r41, %r39, %r40;
; CHECK-NEXT: or.b32 %r42, %r36, %r41;
; CHECK-NEXT: bfe.u32 %r43, %r29, 3, 4;
; CHECK-NEXT: shl.b32 %r44, %r43, 23;
; CHECK-NEXT: or.b32 %r45, %r44, %r31;
; CHECK-NEXT: shl.b32 %r46, %r32, 20;
; CHECK-NEXT: or.b32 %r47, %r45, %r46;
; CHECK-NEXT: add.s32 %r48, %r47, 1006632960;
; CHECK-NEXT: setp.ne.b32 %p6, %r32, 0;
; CHECK-NEXT: selp.b32 %r49, %r42, %r48, %p6;
; CHECK-NEXT: setp.eq.b32 %p7, %r43, 0;
; CHECK-NEXT: selp.b32 %r50, %r49, %r48, %p7;
; CHECK-NEXT: or.b32 %r51, %r43, %r32;
; CHECK-NEXT: setp.eq.b32 %p8, %r51, 0;
; CHECK-NEXT: selp.b32 %r52, %r31, %r50, %p8;
; CHECK-NEXT: setp.eq.b32 %p9, %r32, 7;
; CHECK-NEXT: selp.b32 %r53, 2143289344, %r52, %p9;
; CHECK-NEXT: setp.eq.b32 %p10, %r43, 15;
; CHECK-NEXT: selp.b32 %r54, %r53, %r52, %p10;
; CHECK-NEXT: prmt.b32 %r55, %r1, 0, 0x7771U;
; CHECK-NEXT: shl.b32 %r56, %r55, 24;
; CHECK-NEXT: and.b32 %r57, %r56, -2147483648;
; CHECK-NEXT: and.b32 %r58, %r55, 7;
; CHECK-NEXT: clz.b32 %r59, %r58;
; CHECK-NEXT: sub.s32 %r60, 149, %r59;
; CHECK-NEXT: shl.b32 %r61, %r60, 23;
; CHECK-NEXT: or.b32 %r62, %r57, %r61;
; CHECK-NEXT: sub.s32 %r63, 31, %r59;
; CHECK-NEXT: shl.b32 %r64, %r11, %r63;
; CHECK-NEXT: xor.b32 %r65, %r58, %r64;
; CHECK-NEXT: add.s32 %r66, %r59, -8;
; CHECK-NEXT: shl.b32 %r67, %r65, %r66;
; CHECK-NEXT: or.b32 %r68, %r62, %r67;
; CHECK-NEXT: bfe.u32 %r69, %r55, 3, 4;
; CHECK-NEXT: shl.b32 %r70, %r69, 23;
; CHECK-NEXT: or.b32 %r71, %r70, %r57;
; CHECK-NEXT: shl.b32 %r72, %r58, 20;
; CHECK-NEXT: or.b32 %r73, %r71, %r72;
; CHECK-NEXT: add.s32 %r74, %r73, 1006632960;
; CHECK-NEXT: setp.ne.b32 %p11, %r58, 0;
; CHECK-NEXT: selp.b32 %r75, %r68, %r74, %p11;
; CHECK-NEXT: setp.eq.b32 %p12, %r69, 0;
; CHECK-NEXT: selp.b32 %r76, %r75, %r74, %p12;
; CHECK-NEXT: or.b32 %r77, %r69, %r58;
; CHECK-NEXT: setp.eq.b32 %p13, %r77, 0;
; CHECK-NEXT: selp.b32 %r78, %r57, %r76, %p13;
; CHECK-NEXT: setp.eq.b32 %p14, %r58, 7;
; CHECK-NEXT: selp.b32 %r79, 2143289344, %r78, %p14;
; CHECK-NEXT: setp.eq.b32 %p15, %r69, 15;
; CHECK-NEXT: selp.b32 %r80, %r79, %r78, %p15;
; CHECK-NEXT: prmt.b32 %r81, %r1, 0, 0x7770U;
; CHECK-NEXT: shl.b32 %r82, %r81, 24;
; CHECK-NEXT: and.b32 %r83, %r82, -2147483648;
; CHECK-NEXT: and.b32 %r84, %r81, 7;
; CHECK-NEXT: clz.b32 %r85, %r84;
; CHECK-NEXT: sub.s32 %r86, 149, %r85;
; CHECK-NEXT: shl.b32 %r87, %r86, 23;
; CHECK-NEXT: or.b32 %r88, %r83, %r87;
; CHECK-NEXT: sub.s32 %r89, 31, %r85;
; CHECK-NEXT: shl.b32 %r90, %r11, %r89;
; CHECK-NEXT: xor.b32 %r91, %r84, %r90;
; CHECK-NEXT: add.s32 %r92, %r85, -8;
; CHECK-NEXT: shl.b32 %r93, %r91, %r92;
; CHECK-NEXT: or.b32 %r94, %r88, %r93;
; CHECK-NEXT: bfe.u32 %r95, %r81, 3, 4;
; CHECK-NEXT: shl.b32 %r96, %r95, 23;
; CHECK-NEXT: or.b32 %r97, %r96, %r83;
; CHECK-NEXT: shl.b32 %r98, %r84, 20;
; CHECK-NEXT: or.b32 %r99, %r97, %r98;
; CHECK-NEXT: add.s32 %r100, %r99, 1006632960;
; CHECK-NEXT: setp.ne.b32 %p16, %r84, 0;
; CHECK-NEXT: selp.b32 %r101, %r94, %r100, %p16;
; CHECK-NEXT: setp.eq.b32 %p17, %r95, 0;
; CHECK-NEXT: selp.b32 %r102, %r101, %r100, %p17;
; CHECK-NEXT: or.b32 %r103, %r95, %r84;
; CHECK-NEXT: setp.eq.b32 %p18, %r103, 0;
; CHECK-NEXT: selp.b32 %r104, %r83, %r102, %p18;
; CHECK-NEXT: setp.eq.b32 %p19, %r84, 7;
; CHECK-NEXT: selp.b32 %r105, 2143289344, %r104, %p19;
; CHECK-NEXT: setp.eq.b32 %p20, %r95, 15;
; CHECK-NEXT: selp.b32 %r106, %r105, %r104, %p20;
; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r106, %r80, %r54, %r28};
; CHECK-NEXT: ret;
%r = call <4 x float> @llvm.convert.from.arbitrary.fp.v4f32.v4i8(<4 x i8> %x, metadata !"Float8E4M3FN")
ret <4 x float> %r
}
; <4 x i8> -> <4 x float> Float8E5M2
define <4 x float> @from_f8e5m2_v4f32(<4 x i8> %x) {
; CHECK-LABEL: from_f8e5m2_v4f32(
; CHECK: {
; CHECK-NEXT: .reg .pred %p<21>;
; CHECK-NEXT: .reg .b32 %r<119>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [from_f8e5m2_v4f32_param_0];
; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x7773U;
; CHECK-NEXT: shl.b32 %r3, %r2, 24;
; CHECK-NEXT: and.b32 %r4, %r3, -2147483648;
; CHECK-NEXT: and.b32 %r5, %r2, 3;
; CHECK-NEXT: clz.b32 %r6, %r5;
; CHECK-NEXT: sub.s32 %r7, 142, %r6;
; CHECK-NEXT: shl.b32 %r8, %r7, 23;
; CHECK-NEXT: or.b32 %r9, %r4, %r8;
; CHECK-NEXT: sub.s32 %r10, 31, %r6;
; CHECK-NEXT: mov.b32 %r11, 1;
; CHECK-NEXT: shl.b32 %r12, %r11, %r10;
; CHECK-NEXT: xor.b32 %r13, %r5, %r12;
; CHECK-NEXT: add.s32 %r14, %r6, -8;
; CHECK-NEXT: shl.b32 %r15, %r13, %r14;
; CHECK-NEXT: or.b32 %r16, %r9, %r15;
; CHECK-NEXT: bfe.u32 %r17, %r2, 2, 5;
; CHECK-NEXT: shl.b32 %r18, %r17, 23;
; CHECK-NEXT: or.b32 %r19, %r18, %r4;
; CHECK-NEXT: shl.b32 %r20, %r5, 21;
; CHECK-NEXT: or.b32 %r21, %r19, %r20;
; CHECK-NEXT: add.s32 %r22, %r21, 939524096;
; CHECK-NEXT: setp.ne.b32 %p1, %r5, 0;
; CHECK-NEXT: selp.b32 %r23, %r16, %r22, %p1;
; CHECK-NEXT: setp.eq.b32 %p2, %r17, 0;
; CHECK-NEXT: selp.b32 %r24, %r23, %r22, %p2;
; CHECK-NEXT: or.b32 %r25, %r17, %r5;
; CHECK-NEXT: setp.eq.b32 %p3, %r25, 0;
; CHECK-NEXT: selp.b32 %r26, %r4, %r24, %p3;
; CHECK-NEXT: setp.eq.b32 %p4, %r5, 0;
; CHECK-NEXT: or.b32 %r27, %r4, 2139095040;
; CHECK-NEXT: selp.b32 %r28, %r27, %r26, %p4;
; CHECK-NEXT: setp.eq.b32 %p5, %r17, 31;
; CHECK-NEXT: selp.b32 %r29, %r28, %r26, %p5;
; CHECK-NEXT: selp.b32 %r30, 2143289344, %r29, %p1;
; CHECK-NEXT: selp.b32 %r31, %r30, %r29, %p5;
; CHECK-NEXT: prmt.b32 %r32, %r1, 0, 0x7772U;
; CHECK-NEXT: shl.b32 %r33, %r32, 24;
; CHECK-NEXT: and.b32 %r34, %r33, -2147483648;
; CHECK-NEXT: and.b32 %r35, %r32, 3;
; CHECK-NEXT: clz.b32 %r36, %r35;
; CHECK-NEXT: sub.s32 %r37, 142, %r36;
; CHECK-NEXT: shl.b32 %r38, %r37, 23;
; CHECK-NEXT: or.b32 %r39, %r34, %r38;
; CHECK-NEXT: sub.s32 %r40, 31, %r36;
; CHECK-NEXT: shl.b32 %r41, %r11, %r40;
; CHECK-NEXT: xor.b32 %r42, %r35, %r41;
; CHECK-NEXT: add.s32 %r43, %r36, -8;
; CHECK-NEXT: shl.b32 %r44, %r42, %r43;
; CHECK-NEXT: or.b32 %r45, %r39, %r44;
; CHECK-NEXT: bfe.u32 %r46, %r32, 2, 5;
; CHECK-NEXT: shl.b32 %r47, %r46, 23;
; CHECK-NEXT: or.b32 %r48, %r47, %r34;
; CHECK-NEXT: shl.b32 %r49, %r35, 21;
; CHECK-NEXT: or.b32 %r50, %r48, %r49;
; CHECK-NEXT: add.s32 %r51, %r50, 939524096;
; CHECK-NEXT: setp.ne.b32 %p6, %r35, 0;
; CHECK-NEXT: selp.b32 %r52, %r45, %r51, %p6;
; CHECK-NEXT: setp.eq.b32 %p7, %r46, 0;
; CHECK-NEXT: selp.b32 %r53, %r52, %r51, %p7;
; CHECK-NEXT: or.b32 %r54, %r46, %r35;
; CHECK-NEXT: setp.eq.b32 %p8, %r54, 0;
; CHECK-NEXT: selp.b32 %r55, %r34, %r53, %p8;
; CHECK-NEXT: setp.eq.b32 %p9, %r35, 0;
; CHECK-NEXT: or.b32 %r56, %r34, 2139095040;
; CHECK-NEXT: selp.b32 %r57, %r56, %r55, %p9;
; CHECK-NEXT: setp.eq.b32 %p10, %r46, 31;
; CHECK-NEXT: selp.b32 %r58, %r57, %r55, %p10;
; CHECK-NEXT: selp.b32 %r59, 2143289344, %r58, %p6;
; CHECK-NEXT: selp.b32 %r60, %r59, %r58, %p10;
; CHECK-NEXT: prmt.b32 %r61, %r1, 0, 0x7771U;
; CHECK-NEXT: shl.b32 %r62, %r61, 24;
; CHECK-NEXT: and.b32 %r63, %r62, -2147483648;
; CHECK-NEXT: and.b32 %r64, %r61, 3;
; CHECK-NEXT: clz.b32 %r65, %r64;
; CHECK-NEXT: sub.s32 %r66, 142, %r65;
; CHECK-NEXT: shl.b32 %r67, %r66, 23;
; CHECK-NEXT: or.b32 %r68, %r63, %r67;
; CHECK-NEXT: sub.s32 %r69, 31, %r65;
; CHECK-NEXT: shl.b32 %r70, %r11, %r69;
; CHECK-NEXT: xor.b32 %r71, %r64, %r70;
; CHECK-NEXT: add.s32 %r72, %r65, -8;
; CHECK-NEXT: shl.b32 %r73, %r71, %r72;
; CHECK-NEXT: or.b32 %r74, %r68, %r73;
; CHECK-NEXT: bfe.u32 %r75, %r61, 2, 5;
; CHECK-NEXT: shl.b32 %r76, %r75, 23;
; CHECK-NEXT: or.b32 %r77, %r76, %r63;
; CHECK-NEXT: shl.b32 %r78, %r64, 21;
; CHECK-NEXT: or.b32 %r79, %r77, %r78;
; CHECK-NEXT: add.s32 %r80, %r79, 939524096;
; CHECK-NEXT: setp.ne.b32 %p11, %r64, 0;
; CHECK-NEXT: selp.b32 %r81, %r74, %r80, %p11;
; CHECK-NEXT: setp.eq.b32 %p12, %r75, 0;
; CHECK-NEXT: selp.b32 %r82, %r81, %r80, %p12;
; CHECK-NEXT: or.b32 %r83, %r75, %r64;
; CHECK-NEXT: setp.eq.b32 %p13, %r83, 0;
; CHECK-NEXT: selp.b32 %r84, %r63, %r82, %p13;
; CHECK-NEXT: setp.eq.b32 %p14, %r64, 0;
; CHECK-NEXT: or.b32 %r85, %r63, 2139095040;
; CHECK-NEXT: selp.b32 %r86, %r85, %r84, %p14;
; CHECK-NEXT: setp.eq.b32 %p15, %r75, 31;
; CHECK-NEXT: selp.b32 %r87, %r86, %r84, %p15;
; CHECK-NEXT: selp.b32 %r88, 2143289344, %r87, %p11;
; CHECK-NEXT: selp.b32 %r89, %r88, %r87, %p15;
; CHECK-NEXT: prmt.b32 %r90, %r1, 0, 0x7770U;
; CHECK-NEXT: shl.b32 %r91, %r90, 24;
; CHECK-NEXT: and.b32 %r92, %r91, -2147483648;
; CHECK-NEXT: and.b32 %r93, %r90, 3;
; CHECK-NEXT: clz.b32 %r94, %r93;
; CHECK-NEXT: sub.s32 %r95, 142, %r94;
; CHECK-NEXT: shl.b32 %r96, %r95, 23;
; CHECK-NEXT: or.b32 %r97, %r92, %r96;
; CHECK-NEXT: sub.s32 %r98, 31, %r94;
; CHECK-NEXT: shl.b32 %r99, %r11, %r98;
; CHECK-NEXT: xor.b32 %r100, %r93, %r99;
; CHECK-NEXT: add.s32 %r101, %r94, -8;
; CHECK-NEXT: shl.b32 %r102, %r100, %r101;
; CHECK-NEXT: or.b32 %r103, %r97, %r102;
; CHECK-NEXT: bfe.u32 %r104, %r90, 2, 5;
; CHECK-NEXT: shl.b32 %r105, %r104, 23;
; CHECK-NEXT: or.b32 %r106, %r105, %r92;
; CHECK-NEXT: shl.b32 %r107, %r93, 21;
; CHECK-NEXT: or.b32 %r108, %r106, %r107;
; CHECK-NEXT: add.s32 %r109, %r108, 939524096;
; CHECK-NEXT: setp.ne.b32 %p16, %r93, 0;
; CHECK-NEXT: selp.b32 %r110, %r103, %r109, %p16;
; CHECK-NEXT: setp.eq.b32 %p17, %r104, 0;
; CHECK-NEXT: selp.b32 %r111, %r110, %r109, %p17;
; CHECK-NEXT: or.b32 %r112, %r104, %r93;
; CHECK-NEXT: setp.eq.b32 %p18, %r112, 0;
; CHECK-NEXT: selp.b32 %r113, %r92, %r111, %p18;
; CHECK-NEXT: setp.eq.b32 %p19, %r93, 0;
; CHECK-NEXT: or.b32 %r114, %r92, 2139095040;
; CHECK-NEXT: selp.b32 %r115, %r114, %r113, %p19;
; CHECK-NEXT: setp.eq.b32 %p20, %r104, 31;
; CHECK-NEXT: selp.b32 %r116, %r115, %r113, %p20;
; CHECK-NEXT: selp.b32 %r117, 2143289344, %r116, %p16;
; CHECK-NEXT: selp.b32 %r118, %r117, %r116, %p20;
; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r118, %r89, %r60, %r31};
; CHECK-NEXT: ret;
%r = call <4 x float> @llvm.convert.from.arbitrary.fp.v4f32.v4i8(<4 x i8> %x, metadata !"Float8E5M2")
ret <4 x float> %r
}
; <2 x i8> -> <2 x half> Float8E4M3FN
define <2 x half> @from_f8e4m3fn_v2f16(<2 x i8> %x) {
; CHECK-LABEL: from_f8e4m3fn_v2f16(
; CHECK: {
; CHECK-NEXT: .reg .pred %p<11>;
; CHECK-NEXT: .reg .b16 %rs<52>;
; CHECK-NEXT: .reg .b32 %r<26>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b8 {%rs1, %rs2}, [from_f8e4m3fn_v2f16_param_0];
; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NEXT: and.b32 %r2, %r1, 458759;
; CHECK-NEXT: mov.b32 {%rs3, %rs4}, %r2;
; CHECK-NEXT: cvt.u32.u16 %r3, %rs4;
; CHECK-NEXT: shl.b32 %r4, %r3, 16;
; CHECK-NEXT: clz.b32 %r5, %r4;
; CHECK-NEXT: cvt.u16.u32 %rs5, %r5;
; CHECK-NEXT: sub.s16 %rs6, 15, %rs5;
; CHECK-NEXT: cvt.u32.u16 %r6, %rs6;
; CHECK-NEXT: mov.b16 %rs7, 1;
; CHECK-NEXT: shl.b16 %rs8, %rs7, %r6;
; CHECK-NEXT: cvt.u32.u16 %r7, %rs3;
; CHECK-NEXT: shl.b32 %r8, %r7, 16;
; CHECK-NEXT: clz.b32 %r9, %r8;
; CHECK-NEXT: cvt.u16.u32 %rs9, %r9;
; CHECK-NEXT: sub.s16 %rs10, 15, %rs9;
; CHECK-NEXT: cvt.u32.u16 %r10, %rs10;
; CHECK-NEXT: shl.b16 %rs11, %rs7, %r10;
; CHECK-NEXT: mov.b32 %r11, {%rs11, %rs8};
; CHECK-NEXT: xor.b32 %r12, %r2, %r11;
; CHECK-NEXT: mov.b32 {%rs12, %rs13}, %r12;
; CHECK-NEXT: add.s16 %rs14, %rs5, -5;
; CHECK-NEXT: cvt.u32.u16 %r13, %rs14;
; CHECK-NEXT: shl.b16 %rs15, %rs13, %r13;
; CHECK-NEXT: add.s16 %rs16, %rs9, -5;
; CHECK-NEXT: cvt.u32.u16 %r14, %rs16;
; CHECK-NEXT: shl.b16 %rs17, %rs12, %r14;
; CHECK-NEXT: mov.b32 %r15, {%rs17, %rs15};
; CHECK-NEXT: shl.b16 %rs18, %rs2, 8;
; CHECK-NEXT: and.b16 %rs19, %rs18, -32768;
; CHECK-NEXT: shl.b16 %rs20, %rs1, 8;
; CHECK-NEXT: and.b16 %rs21, %rs20, -32768;
; CHECK-NEXT: mov.b32 %r16, {%rs21, %rs19};
; CHECK-NEXT: sub.s16 %rs22, 21, %rs5;
; CHECK-NEXT: shl.b16 %rs23, %rs22, 10;
; CHECK-NEXT: sub.s16 %rs24, 21, %rs9;
; CHECK-NEXT: shl.b16 %rs25, %rs24, 10;
; CHECK-NEXT: mov.b32 %r17, {%rs25, %rs23};
; CHECK-NEXT: or.b32 %r18, %r16, %r17;
; CHECK-NEXT: or.b32 %r19, %r18, %r15;
; CHECK-NEXT: mov.b32 {%rs26, %rs27}, %r19;
; CHECK-NEXT: shr.u16 %rs28, %rs2, 3;
; CHECK-NEXT: shr.u16 %rs29, %rs1, 3;
; CHECK-NEXT: mov.b32 %r20, {%rs29, %rs28};
; CHECK-NEXT: and.b32 %r21, %r20, 983055;
; CHECK-NEXT: mov.b32 {%rs30, %rs31}, %r21;
; CHECK-NEXT: shl.b16 %rs32, %rs31, 10;
; CHECK-NEXT: add.s16 %rs33, %rs32, 8192;
; CHECK-NEXT: shl.b16 %rs34, %rs30, 10;
; CHECK-NEXT: add.s16 %rs35, %rs34, 8192;
; CHECK-NEXT: mov.b32 %r22, {%rs35, %rs33};
; CHECK-NEXT: or.b32 %r23, %r16, %r22;
; CHECK-NEXT: shl.b16 %rs36, %rs4, 7;
; CHECK-NEXT: shl.b16 %rs37, %rs3, 7;
; CHECK-NEXT: mov.b32 %r24, {%rs37, %rs36};
; CHECK-NEXT: or.b32 %r25, %r23, %r24;
; CHECK-NEXT: mov.b32 {%rs38, %rs39}, %r25;
; CHECK-NEXT: setp.ne.b16 %p1, %rs4, 0;
; CHECK-NEXT: selp.b16 %rs40, %rs27, %rs39, %p1;
; CHECK-NEXT: setp.eq.b16 %p2, %rs31, 0;
; CHECK-NEXT: selp.b16 %rs41, %rs40, %rs39, %p2;
; CHECK-NEXT: or.b16 %rs42, %rs31, %rs4;
; CHECK-NEXT: setp.eq.b16 %p3, %rs42, 0;
; CHECK-NEXT: selp.b16 %rs43, %rs19, %rs41, %p3;
; CHECK-NEXT: setp.eq.b16 %p4, %rs4, 7;
; CHECK-NEXT: selp.b16 %rs44, 32256, %rs43, %p4;
; CHECK-NEXT: setp.eq.b16 %p5, %rs31, 15;
; CHECK-NEXT: selp.b16 %rs45, %rs44, %rs43, %p5;
; CHECK-NEXT: setp.ne.b16 %p6, %rs3, 0;
; CHECK-NEXT: selp.b16 %rs46, %rs26, %rs38, %p6;
; CHECK-NEXT: setp.eq.b16 %p7, %rs30, 0;
; CHECK-NEXT: selp.b16 %rs47, %rs46, %rs38, %p7;
; CHECK-NEXT: or.b16 %rs48, %rs30, %rs3;
; CHECK-NEXT: setp.eq.b16 %p8, %rs48, 0;
; CHECK-NEXT: selp.b16 %rs49, %rs21, %rs47, %p8;
; CHECK-NEXT: setp.eq.b16 %p9, %rs3, 7;
; CHECK-NEXT: selp.b16 %rs50, 32256, %rs49, %p9;
; CHECK-NEXT: setp.eq.b16 %p10, %rs30, 15;
; CHECK-NEXT: selp.b16 %rs51, %rs50, %rs49, %p10;
; CHECK-NEXT: st.param.v2.b16 [func_retval0], {%rs51, %rs45};
; CHECK-NEXT: ret;
%r = call <2 x half> @llvm.convert.from.arbitrary.fp.v2f16.v2i8(<2 x i8> %x, metadata !"Float8E4M3FN")
ret <2 x half> %r
}
define bfloat @from_f8e5m2_bf16(i8 %x) {
; CHECK-LABEL: from_f8e5m2_bf16(
; CHECK: {
; CHECK-NEXT: .reg .pred %p<6>;
; CHECK-NEXT: .reg .b16 %rs<32>;
; CHECK-NEXT: .reg .b32 %r<6>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b8 %rs1, [from_f8e5m2_bf16_param_0];
; CHECK-NEXT: shl.b16 %rs2, %rs1, 8;
; CHECK-NEXT: and.b16 %rs3, %rs2, -32768;
; CHECK-NEXT: and.b16 %rs4, %rs1, 3;
; CHECK-NEXT: cvt.u32.u16 %r1, %rs4;
; CHECK-NEXT: shl.b32 %r2, %r1, 16;
; CHECK-NEXT: clz.b32 %r3, %r2;
; CHECK-NEXT: cvt.u16.u32 %rs5, %r3;
; CHECK-NEXT: sub.s16 %rs6, 126, %rs5;
; CHECK-NEXT: shl.b16 %rs7, %rs6, 7;
; CHECK-NEXT: or.b16 %rs8, %rs3, %rs7;
; CHECK-NEXT: sub.s16 %rs9, 15, %rs5;
; CHECK-NEXT: cvt.u32.u16 %r4, %rs9;
; CHECK-NEXT: mov.b16 %rs10, 1;
; CHECK-NEXT: shl.b16 %rs11, %rs10, %r4;
; CHECK-NEXT: xor.b16 %rs12, %rs4, %rs11;
; CHECK-NEXT: add.s16 %rs13, %rs5, -8;
; CHECK-NEXT: cvt.u32.u16 %r5, %rs13;
; CHECK-NEXT: shl.b16 %rs14, %rs12, %r5;
; CHECK-NEXT: or.b16 %rs15, %rs8, %rs14;
; CHECK-NEXT: shr.u16 %rs16, %rs1, 2;
; CHECK-NEXT: and.b16 %rs17, %rs16, 31;
; CHECK-NEXT: shl.b16 %rs18, %rs17, 7;
; CHECK-NEXT: or.b16 %rs19, %rs18, %rs3;
; CHECK-NEXT: shl.b16 %rs20, %rs4, 5;
; CHECK-NEXT: or.b16 %rs21, %rs19, %rs20;
; CHECK-NEXT: add.s16 %rs22, %rs21, 14336;
; CHECK-NEXT: setp.ne.b16 %p1, %rs4, 0;
; CHECK-NEXT: selp.b16 %rs23, %rs15, %rs22, %p1;
; CHECK-NEXT: setp.eq.b16 %p2, %rs17, 0;
; CHECK-NEXT: selp.b16 %rs24, %rs23, %rs22, %p2;
; CHECK-NEXT: or.b16 %rs25, %rs17, %rs4;
; CHECK-NEXT: setp.eq.b16 %p3, %rs25, 0;
; CHECK-NEXT: selp.b16 %rs26, %rs3, %rs24, %p3;
; CHECK-NEXT: setp.eq.b16 %p4, %rs4, 0;
; CHECK-NEXT: or.b16 %rs27, %rs3, 32640;
; CHECK-NEXT: selp.b16 %rs28, %rs27, %rs26, %p4;
; CHECK-NEXT: setp.eq.b16 %p5, %rs17, 31;
; CHECK-NEXT: selp.b16 %rs29, %rs28, %rs26, %p5;
; CHECK-NEXT: selp.b16 %rs30, 32704, %rs29, %p1;
; CHECK-NEXT: selp.b16 %rs31, %rs30, %rs29, %p5;
; CHECK-NEXT: st.param.b16 [func_retval0], %rs31;
; CHECK-NEXT: ret;
%r = call bfloat @llvm.convert.from.arbitrary.fp.bf16.i8(i8 %x, metadata !"Float8E5M2")
ret bfloat %r
}
define bfloat @from_f8e4m3fn_bf16(i8 %x) {
; CHECK-LABEL: from_f8e4m3fn_bf16(
; CHECK: {
; CHECK-NEXT: .reg .pred %p<6>;
; CHECK-NEXT: .reg .b16 %rs<29>;
; CHECK-NEXT: .reg .b32 %r<6>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b8 %rs1, [from_f8e4m3fn_bf16_param_0];
; CHECK-NEXT: shl.b16 %rs2, %rs1, 8;
; CHECK-NEXT: and.b16 %rs3, %rs2, -32768;
; CHECK-NEXT: and.b16 %rs4, %rs1, 7;
; CHECK-NEXT: cvt.u32.u16 %r1, %rs4;
; CHECK-NEXT: shl.b32 %r2, %r1, 16;
; CHECK-NEXT: clz.b32 %r3, %r2;
; CHECK-NEXT: cvt.u16.u32 %rs5, %r3;
; CHECK-NEXT: sub.s16 %rs6, 133, %rs5;
; CHECK-NEXT: shl.b16 %rs7, %rs6, 7;
; CHECK-NEXT: or.b16 %rs8, %rs3, %rs7;
; CHECK-NEXT: sub.s16 %rs9, 15, %rs5;
; CHECK-NEXT: cvt.u32.u16 %r4, %rs9;
; CHECK-NEXT: mov.b16 %rs10, 1;
; CHECK-NEXT: shl.b16 %rs11, %rs10, %r4;
; CHECK-NEXT: xor.b16 %rs12, %rs4, %rs11;
; CHECK-NEXT: add.s16 %rs13, %rs5, -8;
; CHECK-NEXT: cvt.u32.u16 %r5, %rs13;
; CHECK-NEXT: shl.b16 %rs14, %rs12, %r5;
; CHECK-NEXT: or.b16 %rs15, %rs8, %rs14;
; CHECK-NEXT: shr.u16 %rs16, %rs1, 3;
; CHECK-NEXT: and.b16 %rs17, %rs16, 15;
; CHECK-NEXT: shl.b16 %rs18, %rs17, 7;
; CHECK-NEXT: or.b16 %rs19, %rs18, %rs3;
; CHECK-NEXT: shl.b16 %rs20, %rs4, 4;
; CHECK-NEXT: or.b16 %rs21, %rs19, %rs20;
; CHECK-NEXT: add.s16 %rs22, %rs21, 15360;
; CHECK-NEXT: setp.ne.b16 %p1, %rs4, 0;
; CHECK-NEXT: selp.b16 %rs23, %rs15, %rs22, %p1;
; CHECK-NEXT: setp.eq.b16 %p2, %rs17, 0;
; CHECK-NEXT: selp.b16 %rs24, %rs23, %rs22, %p2;
; CHECK-NEXT: or.b16 %rs25, %rs17, %rs4;
; CHECK-NEXT: setp.eq.b16 %p3, %rs25, 0;
; CHECK-NEXT: selp.b16 %rs26, %rs3, %rs24, %p3;
; CHECK-NEXT: setp.eq.b16 %p4, %rs4, 7;
; CHECK-NEXT: selp.b16 %rs27, 32704, %rs26, %p4;
; CHECK-NEXT: setp.eq.b16 %p5, %rs17, 15;
; CHECK-NEXT: selp.b16 %rs28, %rs27, %rs26, %p5;
; CHECK-NEXT: st.param.b16 [func_retval0], %rs28;
; CHECK-NEXT: ret;
%r = call bfloat @llvm.convert.from.arbitrary.fp.bf16.i8(i8 %x, metadata !"Float8E4M3FN")
ret bfloat %r
}
define <2 x bfloat> @from_f8e5m2_v2bf16(<2 x i8> %x) {
; CHECK-LABEL: from_f8e5m2_v2bf16(
; CHECK: {
; CHECK-NEXT: .reg .pred %p<11>;
; CHECK-NEXT: .reg .b16 %rs<58>;
; CHECK-NEXT: .reg .b32 %r<27>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b8 {%rs1, %rs2}, [from_f8e5m2_v2bf16_param_0];
; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NEXT: and.b32 %r2, %r1, 196611;
; CHECK-NEXT: mov.b32 {%rs3, %rs4}, %r2;
; CHECK-NEXT: cvt.u32.u16 %r3, %rs4;
; CHECK-NEXT: shl.b32 %r4, %r3, 16;
; CHECK-NEXT: clz.b32 %r5, %r4;
; CHECK-NEXT: cvt.u16.u32 %rs5, %r5;
; CHECK-NEXT: sub.s16 %rs6, 15, %rs5;
; CHECK-NEXT: cvt.u32.u16 %r6, %rs6;
; CHECK-NEXT: mov.b16 %rs7, 1;
; CHECK-NEXT: shl.b16 %rs8, %rs7, %r6;
; CHECK-NEXT: cvt.u32.u16 %r7, %rs3;
; CHECK-NEXT: shl.b32 %r8, %r7, 16;
; CHECK-NEXT: clz.b32 %r9, %r8;
; CHECK-NEXT: cvt.u16.u32 %rs9, %r9;
; CHECK-NEXT: sub.s16 %rs10, 15, %rs9;
; CHECK-NEXT: cvt.u32.u16 %r10, %rs10;
; CHECK-NEXT: shl.b16 %rs11, %rs7, %r10;
; CHECK-NEXT: mov.b32 %r11, {%rs11, %rs8};
; CHECK-NEXT: xor.b32 %r12, %r2, %r11;
; CHECK-NEXT: mov.b32 {%rs12, %rs13}, %r12;
; CHECK-NEXT: add.s16 %rs14, %rs5, -8;
; CHECK-NEXT: cvt.u32.u16 %r13, %rs14;
; CHECK-NEXT: shl.b16 %rs15, %rs13, %r13;
; CHECK-NEXT: add.s16 %rs16, %rs9, -8;
; CHECK-NEXT: cvt.u32.u16 %r14, %rs16;
; CHECK-NEXT: shl.b16 %rs17, %rs12, %r14;
; CHECK-NEXT: mov.b32 %r15, {%rs17, %rs15};
; CHECK-NEXT: shl.b16 %rs18, %rs2, 8;
; CHECK-NEXT: and.b16 %rs19, %rs18, -32768;
; CHECK-NEXT: shl.b16 %rs20, %rs1, 8;
; CHECK-NEXT: and.b16 %rs21, %rs20, -32768;
; CHECK-NEXT: mov.b32 %r16, {%rs21, %rs19};
; CHECK-NEXT: sub.s16 %rs22, 126, %rs5;
; CHECK-NEXT: shl.b16 %rs23, %rs22, 7;
; CHECK-NEXT: sub.s16 %rs24, 126, %rs9;
; CHECK-NEXT: shl.b16 %rs25, %rs24, 7;
; CHECK-NEXT: mov.b32 %r17, {%rs25, %rs23};
; CHECK-NEXT: or.b32 %r18, %r16, %r17;
; CHECK-NEXT: or.b32 %r19, %r18, %r15;
; CHECK-NEXT: mov.b32 {%rs26, %rs27}, %r19;
; CHECK-NEXT: shr.u16 %rs28, %rs2, 2;
; CHECK-NEXT: shr.u16 %rs29, %rs1, 2;
; CHECK-NEXT: mov.b32 %r20, {%rs29, %rs28};
; CHECK-NEXT: and.b32 %r21, %r20, 2031647;
; CHECK-NEXT: mov.b32 {%rs30, %rs31}, %r21;
; CHECK-NEXT: shl.b16 %rs32, %rs31, 7;
; CHECK-NEXT: add.s16 %rs33, %rs32, 14336;
; CHECK-NEXT: shl.b16 %rs34, %rs30, 7;
; CHECK-NEXT: add.s16 %rs35, %rs34, 14336;
; CHECK-NEXT: mov.b32 %r22, {%rs35, %rs33};
; CHECK-NEXT: or.b32 %r23, %r16, %r22;
; CHECK-NEXT: shl.b16 %rs36, %rs4, 5;
; CHECK-NEXT: shl.b16 %rs37, %rs3, 5;
; CHECK-NEXT: mov.b32 %r24, {%rs37, %rs36};
; CHECK-NEXT: or.b32 %r25, %r23, %r24;
; CHECK-NEXT: mov.b32 {%rs38, %rs39}, %r25;
; CHECK-NEXT: setp.ne.b16 %p1, %rs4, 0;
; CHECK-NEXT: selp.b16 %rs40, %rs27, %rs39, %p1;
; CHECK-NEXT: setp.eq.b16 %p2, %rs31, 0;
; CHECK-NEXT: selp.b16 %rs41, %rs40, %rs39, %p2;
; CHECK-NEXT: or.b16 %rs42, %rs31, %rs4;
; CHECK-NEXT: setp.eq.b16 %p3, %rs42, 0;
; CHECK-NEXT: selp.b16 %rs43, %rs19, %rs41, %p3;
; CHECK-NEXT: or.b32 %r26, %r16, 2139127680;
; CHECK-NEXT: mov.b32 {%rs44, %rs45}, %r26;
; CHECK-NEXT: setp.eq.b16 %p4, %rs4, 0;
; CHECK-NEXT: selp.b16 %rs46, %rs45, %rs43, %p4;
; CHECK-NEXT: setp.eq.b16 %p5, %rs31, 31;
; CHECK-NEXT: selp.b16 %rs47, %rs46, %rs43, %p5;
; CHECK-NEXT: selp.b16 %rs48, 32704, %rs47, %p1;
; CHECK-NEXT: selp.b16 %rs49, %rs48, %rs47, %p5;
; CHECK-NEXT: setp.ne.b16 %p6, %rs3, 0;
; CHECK-NEXT: selp.b16 %rs50, %rs26, %rs38, %p6;
; CHECK-NEXT: setp.eq.b16 %p7, %rs30, 0;
; CHECK-NEXT: selp.b16 %rs51, %rs50, %rs38, %p7;
; CHECK-NEXT: or.b16 %rs52, %rs30, %rs3;
; CHECK-NEXT: setp.eq.b16 %p8, %rs52, 0;
; CHECK-NEXT: selp.b16 %rs53, %rs21, %rs51, %p8;
; CHECK-NEXT: setp.eq.b16 %p9, %rs3, 0;
; CHECK-NEXT: selp.b16 %rs54, %rs44, %rs53, %p9;
; CHECK-NEXT: setp.eq.b16 %p10, %rs30, 31;
; CHECK-NEXT: selp.b16 %rs55, %rs54, %rs53, %p10;
; CHECK-NEXT: selp.b16 %rs56, 32704, %rs55, %p6;
; CHECK-NEXT: selp.b16 %rs57, %rs56, %rs55, %p10;
; CHECK-NEXT: st.param.v2.b16 [func_retval0], {%rs57, %rs49};
; CHECK-NEXT: ret;
%r = call <2 x bfloat> @llvm.convert.from.arbitrary.fp.v2bf16.v2i8(<2 x i8> %x, metadata !"Float8E5M2")
ret <2 x bfloat> %r
}