; 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 <4 x float> @llvm.convert.from.arbitrary.fp.v4f32.v4i4(<4 x i4>, metadata)

declare half @llvm.convert.from.arbitrary.fp.f16.i8(i8, metadata)
declare double @llvm.convert.from.arbitrary.fp.f64.i8(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
}