From 665e88008b091be21b0734f4b5a0b3e0417734e8 Mon Sep 17 00:00:00 2001 From: "Chaudhary, Jatin Jaikishan" Date: Thu, 10 Apr 2025 17:22:15 +0100 Subject: [PATCH] SWDEV-461087 - fp4/fp6/fp8 ocp headers (#41) This now has host conversions too, which is directly from Christopher's work on fcbx. Signed-off-by: Christopher M. Riedl * add const to func parameter * do not depend on builtins, use gfx950 detection [ROCm/clr commit: 628777b73d27f6d2c1abd005e8c5ed9737ec76c6] --- .../include/hip/amd_detail/amd_hip_ocp_fp.hpp | 2886 +++++++++++++++++ .../hip/amd_detail/amd_hip_ocp_fp_cxx.hpp | 1056 ++++++ .../hip/amd_detail/amd_hip_ocp_host.hpp | 961 ++++++ .../hip/amd_detail/amd_hip_ocp_types.h | 66 + 4 files changed, 4969 insertions(+) create mode 100644 projects/clr/hipamd/include/hip/amd_detail/amd_hip_ocp_fp.hpp create mode 100644 projects/clr/hipamd/include/hip/amd_detail/amd_hip_ocp_fp_cxx.hpp create mode 100644 projects/clr/hipamd/include/hip/amd_detail/amd_hip_ocp_host.hpp create mode 100644 projects/clr/hipamd/include/hip/amd_detail/amd_hip_ocp_types.h diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_ocp_fp.hpp b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_ocp_fp.hpp new file mode 100644 index 0000000000..757da7ea26 --- /dev/null +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_ocp_fp.hpp @@ -0,0 +1,2886 @@ +#pragma once + +#include + +#if !defined(__HIPCC_RTC__) +#include +#include +#include +#include +#include +#include +#include + +static_assert(sizeof(uint8_t) * CHAR_BIT == 8); +static_assert(sizeof(uint16_t) * CHAR_BIT == 16); +static_assert(sizeof(uint32_t) * CHAR_BIT == 32); +static_assert(sizeof(uint64_t) * CHAR_BIT == 64); +#endif // !defined(__HIPCC_RTC__) + +#include // Host Conversion + +// HW Detection +#if defined(__gfx950__) +#define HIP_ENABLE_GFX950_OCP_BUILTINS 1 +#else +#define HIP_ENABLE_GFX950_OCP_BUILTINS 0 +#endif +#if !defined(__gfx950__) +#define HIP_ENABLE_HOST_OCP_CONVERSIONS 1 +#else +#define HIP_ENABLE_HOST_OCP_CONVERSIONS 0 +#endif + +/** FP8 format + * +---------------+--------------------+----------------------+ + * | Attribute | E4M3 | E5M2 | + * +---------------+--------------------+----------------------+ + * | Exponent Bias | 7 | 15 | + * | Inf | N/A | S-11111-00 | + * | NAN | S-1111-111 | S-11111-** | + * | Zero | S-0000-000 | S-00000-00 | + * | Max Normal | S-1111-110: +-448 | S-11110-11: +-57,344 | + * | Min Normal | S-0001-000: +-2^-6 | S-00001-00: +-2^-14 | + * | Max Subnorm | S-0000-111 | S-00000-11 | + * | Min Subnorm | S-0000-001 | S-00000-01 | + * +---------------+--------------------+----------------------+ + */ +enum __amd_fp8_interpretation_t { + __AMD_OCP_E4M3 = 0, /* FP8 */ + __AMD_OCP_E5M2 = 1, /* BF8 */ +}; + +/** FP6 format + * +---------------+------------------------------------+-------------------------------------+ + * | Attribute | E2M3 | E3M2 | + * +---------------+------------------------------------+-------------------------------------+ + * | Exponent bias | 1 | 3 | + * | Infinities | N/A | N/A | + * | NaN | N/A | N/A | + * | Zeros | S-00-000 | S-000-00 | + * | Max normal | S-11-111 = +-2^2 × 1.875 = +-7.5 | S-111-11 = +-2^4 × 1.75 = +-28.0 | + * | Min normal | S-01-000 = +-2^0 × 1.0 = +-1.0 | S-001-00 = +-2^-2 × 1.0 = +-0.25 | + * | Max subnorm | S-00-111 = +-2^0 × 0.875 = +-0.875 | S-000-11 = +-2^-2 × 0.75 = +-0.1875 | + * | Min subnorm | S-00-001 = +-2^0 × 0.125 = +-0.125 | S-000-01 = +-2^-2 × 0.25 = +-0.0625 | + * +---------------+------------------------------------+-------------------------------------+ + */ +enum __amd_fp6_interpretation_t { + __AMD_OCP_E2M3 = 0, /* FP6 */ + __AMD_OCP_E3M2 = 1, /* BF6 */ +}; + +/** FP4 format + * +---------------+------------------------------+ + * | Attribute | E2M1 | + * +---------------+------------------------------+ + * | Exponent bias | 1 | + * | Infinities | N/A | + * | NaN | N/A | + * | Zeros | S-00-0 | + * | Max normal | S-11-1 = +-2^2 × 1.5 = +-6.0 | + * | Min normal | S-01-0 = +-2^0 × 1.0 = +-1.0 | + * | Max subnorm | S-00-1 = +-2^0 × 0.5 = +-0.5 | + * | Min subnorm | S-00-1 = +-2^0 × 0.5 = +-0.5 | + * +---------------+------------------------------+ + */ +enum __amd_fp4_interpretation_t { + __AMD_OCP_E2M1 = 0, /* FP4 */ +}; + +/** + * @brief Create fp8x2 from two fp8 numbers + * + * @param x + * @param y + * @return __amd_fp8x2_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp8x2_storage_t +__amd_create_fp8x2(const __amd_fp8_storage_t x, const __amd_fp8_storage_t y) { + __amd_fp8x2_storage_t ret = 0; + ret = x | (y << 8); + return ret; +} + +/** + * @brief Create fp4x2 from two fp4 numbers + * + * @param x + * @param y + * @return __amd_fp4x2_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp4x2_storage_t __amd_create_fp4x2(const uint8_t x, + const uint8_t y) { + __amd_fp4x2_storage_t ret = 0; + ret = x | (y << 4); + return ret; +} + +/** + * @brief Create fp4x8 from fp4x2 + * + * @param x + * @param y + * @param z + * @param w + * @return __amd_fp4x8_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp4x8_storage_t +__amd_create_fp4x8(const __amd_fp4x2_storage_t x, const __amd_fp4x2_storage_t y, + const __amd_fp4x2_storage_t z, const __amd_fp4x2_storage_t w) { + static_assert(sizeof(__amd_fp4x2_storage_t[4]) == sizeof(__amd_fp4x8_storage_t)); + union { + __amd_fp4x2_storage_t fp4x2[4]; + __amd_fp4x8_storage_t fp4x8; + } u{{x, y, z, w}}; + return u.fp4x8; +} + +/** + * @brief Create fp8x2 from two fp8x2 numbers + * + * @param x + * @param y + * @return __amd_fp8x8_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp8x8_storage_t +__amd_create_fp8x8(const __amd_fp8x2_storage_t x, const __amd_fp8x2_storage_t y, + const __amd_fp8x2_storage_t z, const __amd_fp8x2_storage_t w) { + static_assert(sizeof(__amd_fp8x2_storage_t[4]) == sizeof(__amd_fp8x8_storage_t)); + union { + __amd_fp8x2_storage_t fp8x2[4]; + __amd_fp8x8_storage_t fp8x8; + } u{{x, y, z, w}}; + return u.fp8x8; +} + +/** + * @brief Get fp8 number from an fp8x2 + * + * @param x + * @param index + * @return __amd_fp8_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp8_storage_t __amd_extract_fp8(const __amd_fp8x2_storage_t x, + const size_t index) { + static_assert(sizeof(__amd_fp8_storage_t[2]) == sizeof(__amd_fp8x2_storage_t)); + union { + __amd_fp8x2_storage_t fp8x2; + __amd_fp8_storage_t fp8[2]; + } u{x}; + if (index == 0) return u.fp8[0]; + return u.fp8[1]; +} + +/** + * @brief Get fp8x2 from fp8x4 + * + * @param x + * @param index + * @return __amd_fp8x2_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp8x2_storage_t +__amd_extract_fp8x2(const __amd_fp8x8_storage_t x, const size_t index) { + static_assert(sizeof(__amd_fp8x8_storage_t) == sizeof(__amd_fp8x2_storage_t[4])); + union { + __amd_fp8x8_storage_t fp8x8; + __amd_fp8x2_storage_t fp8x2[4]; + } u{x}; + return u.fp8x2[index]; +} + +/** + * @brief Extract a fp4 from fp4x2 + * + * @param x __amd_fp4x2_storage_t type + * @param index 0 or 1 + * @return uint8_t populated with 4 bits of fp4 number + */ +__OCP_FP_HOST_DEVICE_STATIC__ uint8_t __amd_extract_fp4(const __amd_fp4x2_storage_t x, + const size_t index) { + if (index == 0) return (x & 0xFu); + return (x >> 4); +} + +/** + * @brief extract fp4x2 from fp4x4 + * + * @param x __amd_fp4x8_storage_t type + * @param index 0, 1, 2 or 3 index + * @return __amd_fp4x2_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp4x2_storage_t +__amd_extract_fp4x2(const __amd_fp4x8_storage_t x, const size_t index) { + return (x >> (index * 8 /* char bits*/)) & 0xFFu; +} + +/** + * @brief Convert scale E8M0 to float scale which is used by gfx950 + * + * @param scale + * @return float embed with E8M0 in its exponent + */ +__OCP_FP_HOST_DEVICE_STATIC__ float __amd_scale_to_float(const __amd_scale_t scale_exp) { + constexpr int8_t OCP_SCALE_EXP_NAN = -128; + const uint32_t SCALE_EXP_BIAS = 127; // OCP MX E8M0 "scale" bias + + // On gfx950 the "scale" operand is encoded in the exponent bits of + // an IEEE-754 float - always in the form: 1.0 * 2**scale. + const size_t SCALE_EXP_SHIFT = 23; // IEEE-754 "float" mantissa bits + + uint32_t s; + if (scale_exp == OCP_SCALE_EXP_NAN) + s = 0xff; + else + s = ((uint32_t)scale_exp + SCALE_EXP_BIAS) & 0xffu; + + return fcbx::F32(s << SCALE_EXP_SHIFT); +} + +/** + * @brief Convert FP8 to float + * + * @param val input fp8 val + * @param interpret + * @return float + */ +__OCP_FP_HOST_DEVICE_STATIC__ float __amd_cvt_fp8_to_float( + const __amd_fp8_storage_t val, const __amd_fp8_interpretation_t interpret) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + if (interpret == __AMD_OCP_E4M3) { + return __builtin_amdgcn_cvt_f32_fp8(val, 0); + } else { + return __builtin_amdgcn_cvt_f32_bf8(val, 0); + } +#else // Host + using namespace fcbx; + if (interpret == __AMD_OCP_E4M3) { + return to_float(static_cast(val), 0); + } else { + return to_float(static_cast(val), 0); + } +#endif +} + +/** + * @brief Convert float to FP8 with stochastic rounding. + * + * @param val input float value. + * @param interpret + * @param seed + * @return __amd_fp8_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp8_storage_t __amd_cvt_float_to_fp8_sr( + const float val, const __amd_fp8_interpretation_t interpret, const unsigned int seed) { + static_assert(sizeof(float) == sizeof(__amd_fp8_storage_t[4])); + static_assert(sizeof(unsigned int) == sizeof(uint32_t)); + union { + unsigned int ui32; + uint32_t ui32t; + float f32; + __amd_fp8_storage_t fp8[4]; + } u{0}; +#if HIP_ENABLE_GFX950_OCP_BUILTINS + if (interpret == __AMD_OCP_E4M3) { + u.ui32 = __builtin_amdgcn_cvt_sr_fp8_f32(val, seed, u.ui32, 0); + } else { + u.ui32 = __builtin_amdgcn_cvt_sr_bf8_f32(val, seed, u.ui32, 0); + } +#else + using namespace fcbx; + if (interpret == __AMD_OCP_E4M3) { + u.ui32t = from_float_sr(val, seed, 0 /*scale*/); + } else { + u.ui32t = from_float_sr(val, seed, 0 /*scale*/); + } +#endif + return u.fp8[0]; +} + + +/** + * @brief Convert fp8 to float with scale. + * + * @param in fp8 number + * @param interpret + * @param scale + * @return float + */ +__OCP_FP_HOST_DEVICE_STATIC__ float __amd_cvt_fp8_to_float_scale( + const __amd_fp8_storage_t val, const __amd_fp8_interpretation_t interpret, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + return interpret == __AMD_OCP_E4M3 + ? __builtin_amdgcn_cvt_scalef32_f32_fp8(val, __amd_scale_to_float(scale), 0) + : __builtin_amdgcn_cvt_scalef32_f32_bf8(val, __amd_scale_to_float(scale), 0); +#else + using namespace fcbx; + return interpret == __AMD_OCP_E4M3 + ? to_float(static_cast(val), scale) + : to_float(static_cast(val), scale); +#endif +} + +/** + * @brief Convert float to fp8 with stochastic rounding and scale. + * + * @param val + * @param interpret + * @param seed + * @param scale + * @return __amd_fp8_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp8_storage_t +__amd_cvt_float_to_fp8_sr_scale(const float val, const __amd_fp8_interpretation_t interpret, + const unsigned int seed, const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(__amd_fp8_storage_t[4]) == sizeof(unsigned int)); + union u { + unsigned int ui32; + __amd_fp8_storage_t fp8[4]; + } u{0}; + if (interpret == __AMD_OCP_E4M3) { + u.ui32 = + __builtin_amdgcn_cvt_scalef32_sr_fp8_f32(u.ui32, val, seed, __amd_scale_to_float(scale), 0); + } else { + u.ui32 = + __builtin_amdgcn_cvt_scalef32_sr_bf8_f32(u.ui32, val, seed, __amd_scale_to_float(scale), 0); + } + return u.fp8[0]; +#else + static_assert(sizeof(uint32_t) == sizeof(unsigned int)); + union u { + uint32_t ui32t; + __amd_fp8_storage_t fp8[4]; + } u{0}; + using namespace fcbx; + u.ui32t = interpret == __AMD_OCP_E4M3 + ? from_float_sr(val, seed, scale) + : from_float_sr(val, seed, scale); + return u.fp8[0]; +#endif +} + +/** + * @brief Convert packed fp8x2 to floatx2. + * + * @param val input fp8x2 value + * @param interpret + * @return __amd_floatx2_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_floatx2_storage_t __amd_cvt_fp8x2_to_floatx2( + const __amd_fp8x2_storage_t val, const __amd_fp8_interpretation_t interpret) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + if (interpret == __AMD_OCP_E4M3) { + return __builtin_amdgcn_cvt_pk_f32_fp8(val, false); + } else { + return __builtin_amdgcn_cvt_pk_f32_bf8(val, false); + } +#else + using namespace fcbx; + __amd_floatx2_storage_t ret; + if (interpret == __AMD_OCP_E4M3) { + ret[0] = + to_float(static_cast(__amd_extract_fp8(val, 0)), 0); + ret[1] = + to_float(static_cast(__amd_extract_fp8(val, 1)), 0); + } else { + ret[0] = + to_float(static_cast(__amd_extract_fp8(val, 0)), 0); + ret[1] = + to_float(static_cast(__amd_extract_fp8(val, 1)), 0); + } + return ret; +#endif +} + +/** + * @brief Convert floatx2 to fp8x2. + * + * @param val Input floatx2 value. + * @param interpret + * @return __amd_fp8x2_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp8x2_storage_t __amd_cvt_floatx2_to_fp8x2( + const __amd_floatx2_storage_t val, const __amd_fp8_interpretation_t interpret) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(__amd_fp8x2_storage_t[2]) == sizeof(int)); + union { + int i32; + __amd_fp8x2_storage_t fp8x2[2]; + } u{0}; + if (interpret == __AMD_OCP_E4M3) { + u.i32 = __builtin_amdgcn_cvt_pk_fp8_f32(val[0], val[1], u.i32, false); + } else { + u.i32 = __builtin_amdgcn_cvt_pk_bf8_f32(val[0], val[1], u.i32, false); + } + return u.fp8x2[0]; +#else + using namespace fcbx; + __amd_fp8_storage_t l, r; + if (interpret == __AMD_OCP_E4M3) { + l = from_float(val[0], 0 /*scale*/); + r = from_float(val[1], 0 /*scale*/); + } else { + l = from_float(val[0], 0 /*scale*/); + r = from_float(val[1], 0 /*scale*/); + } + return __amd_create_fp8x2(l, r); +#endif +} + +/** + * @brief Convert packed floatx2 to fp4x2 with stochastic rounding and scale. + * + * @param val + * @param seed + * @param scale + * @return __amd_fp4x2_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp4x2_storage_t __amd_cvt_floatx2_to_fp4x2_sr_scale( + const __amd_floatx2_storage_t val, const __amd_fp4_interpretation_t, const unsigned int seed, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(__amd_fp4x2_storage_t[4]) == sizeof(unsigned int)); + union { + unsigned int ui32; + __amd_fp4x2_storage_t fp4x2[4]; + } u{0}; + u.ui32 = __builtin_amdgcn_cvt_scalef32_sr_pk_fp4_f32(u.ui32, val, seed, + __amd_scale_to_float(scale), 1); + return u.fp4x2[1]; +#else + static_assert(sizeof(__amd_fp4x2_storage_t[4]) == sizeof(uint32_t)); + union u { + uint32_t ui32t; + __amd_fp4x2_storage_t fp4x2[4]; + } u{0}; + using namespace fcbx; + auto l = from_float_sr(val[0], seed, scale); + auto r = from_float_sr(val[1], seed, scale); + r <<= 4; + l |= r; + u.ui32t = l; + return u.fp4x2[0]; +#endif +} + +/** + * @brief Convert packed fp4x2 to floatx2. This is wrapper for gfx950 builtin. + * + * @param in + * @param scale + * @return __amd_floatx2_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_floatx2_storage_t __amd_cvt_fp4x2_to_floatx2_scale( + const __amd_fp4x2_storage_t val, const __amd_fp4_interpretation_t, const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + return __builtin_amdgcn_cvt_scalef32_pk_f32_fp4(val, __amd_scale_to_float(scale), 0); +#else + using namespace fcbx; + __amd_floatx2_storage_t ret{to_float(val & 0xFu, scale), + to_float(val >> 4, scale)}; + return ret; +#endif +} + +/** + * @brief Convert packed floatx2 to fp4x2. + * + * @param in + * @param scale + * @return __amd_fp4x2_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp4x2_storage_t +__amd_cvt_floatx2_to_fp4x2_scale(const __amd_floatx2_storage_t val, + const __amd_fp4_interpretation_t, const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(unsigned int) == sizeof(__amd_fp4x2_storage_t[4])); + union { + unsigned int ui32; + __amd_fp4x2_storage_t fp4x2[4]; + } u{0}; + u.ui32 = __builtin_amdgcn_cvt_scalef32_pk_fp4_f32(u.ui32, val[0], val[1], + __amd_scale_to_float(scale), 0); + return u.fp4x2[0]; +#else + using namespace fcbx; + auto l = from_float(val[1], scale); + auto r = from_float(val[0], scale); + __amd_fp4x2_storage_t ret(l << 4 | r); + return ret; +#endif +} + +/** + * @brief Convert packed fp8x2 to floatx2. This is direct mapping of gfx950. + * + * @param in input fp8x2 + * @param interpret + * @param scale + * @return __amd_floatx2_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_floatx2_storage_t __amd_cvt_fp8x2_to_floatx2_scale( + const __amd_fp8x2_storage_t val, const __amd_fp8_interpretation_t interpret, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + return interpret == __AMD_OCP_E4M3 + ? __builtin_amdgcn_cvt_scalef32_pk_f32_fp8(val, __amd_scale_to_float(scale), false) + : __builtin_amdgcn_cvt_scalef32_pk_f32_bf8(val, __amd_scale_to_float(scale), false); +#else + using namespace fcbx; + __amd_floatx2_storage_t ret; + if (interpret == __AMD_OCP_E4M3) { + ret[0] = to_float(val >> 8, scale); + ret[1] = to_float((val << 8) >> 8, scale); + } else { + ret[0] = to_float(val >> 8, scale); + ret[1] = to_float((val << 8) >> 8, scale); + } + return ret; +#endif +} + +/** + * @brief Convert packed floatx2 to fp8x2 with scale. + * + * @param in + * @param interpret + * @param scale + * @return __amd_fp8x2_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp8x2_storage_t __amd_cvt_floatx2_to_fp8x2_scale( + const __amd_floatx2_storage_t val, const __amd_fp8_interpretation_t interpret, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(__amd_shortx2_storage_t) == sizeof(__amd_fp8x2_storage_t[2])); + union { + __amd_shortx2_storage_t shortx2; + __amd_fp8x2_storage_t fp8x2[2]; + } u{0}; + u.shortx2 = interpret == __AMD_OCP_E4M3 + ? __builtin_amdgcn_cvt_scalef32_pk_fp8_f32(u.shortx2, val[0], val[1], + __amd_scale_to_float(scale), false) + : __builtin_amdgcn_cvt_scalef32_pk_bf8_f32(u.shortx2, val[0], val[1], + __amd_scale_to_float(scale), false); + return u.fp8x2[0]; +#else + using namespace fcbx; + uint8_t l, r; + if (interpret == __AMD_OCP_E4M3) { + l = from_float(val[0], scale); + r = from_float(val[1], scale); + } else { + l = from_float(val[0], scale); + r = from_float(val[1], scale); + } + __amd_fp8x2_storage_t ret(l << 8 | r); + return ret; +#endif +} + +/** + * @brief Convert packed bf16x32 to fp6x32 with scale. + * + * @param in + * @param interpret + * @param scale + * @return __amd_fp6x32_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp6x32_storage_t __amd_cvt_bf16x32_to_fp6x32_scale( + const __amd_bf16x32_storage_t in, const __amd_fp6_interpretation_t interpret, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + if (interpret == __AMD_OCP_E2M3) { + return __builtin_amdgcn_cvt_scalef32_pk32_fp6_bf16(in, __amd_scale_to_float(scale)); + } + return __builtin_amdgcn_cvt_scalef32_pk32_bf6_bf16(in, __amd_scale_to_float(scale)); +#else + if (interpret == __AMD_OCP_E2M3) { + return fcbx::fp6_cvt_packedx32<__amd_bf16x32_storage_t, __amd_fp6x32_storage_t, + __amd_bf16_storage_t, fcbx::Encoding::E8M7, + fcbx::Encoding::E2M3>(in, scale); + } else { + return fcbx::fp6_cvt_packedx32<__amd_bf16x32_storage_t, __amd_fp6x32_storage_t, + __amd_bf16_storage_t, fcbx::Encoding::E8M7, + fcbx::Encoding::E3M2>(in, scale); + } +#endif +} + +/** + * @brief Convert packed fp16x32 to fp6x32 with scale. + * + * @param in + * @param interpret + * @param scale + * @return __amd_fp6x32_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp6x32_storage_t __amd_cvt_fp16x32_to_fp6x32_scale( + const __amd_fp16x32_storage_t in, const __amd_fp6_interpretation_t interpret, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + if (interpret == __AMD_OCP_E2M3) { + return __builtin_amdgcn_cvt_scalef32_pk32_fp6_f16(in, __amd_scale_to_float(scale)); + } + return __builtin_amdgcn_cvt_scalef32_pk32_bf6_f16(in, __amd_scale_to_float(scale)); +#else + if (interpret == __AMD_OCP_E2M3) { + return fcbx::fp6_cvt_packedx32<__amd_fp16x32_storage_t, __amd_fp6x32_storage_t, + __amd_fp16_storage_t, fcbx::Encoding::E5M10, + fcbx::Encoding::E2M3>(in, scale); + } else { + return fcbx::fp6_cvt_packedx32<__amd_fp16x32_storage_t, __amd_fp6x32_storage_t, + __amd_fp16_storage_t, fcbx::Encoding::E5M10, + fcbx::Encoding::E3M2>(in, scale); + } +#endif +} + +/** + * @brief Convert fp8x2 to fp16x2. This is direct mapping of gfx950 builtin. + * + * @param val + * @param interpret + * @param scale + * @return __amd_fp16x2_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp16x2_storage_t __amd_cvt_fp8x2_to_fp16x2_scale( + const __amd_fp8x2_storage_t val, const __amd_fp8_interpretation_t interpret, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(unsigned int) == sizeof(__amd_fp8x2_storage_t[2])); + union { + __amd_fp8x2_storage_t fp8x2[2]; + unsigned int ui32; + } u; + u.fp8x2[0] = val; + return interpret == __AMD_OCP_E4M3 + ? __builtin_amdgcn_cvt_scalef32_pk_f16_fp8(u.ui32, __amd_scale_to_float(scale), false) + : __builtin_amdgcn_cvt_scalef32_pk_f16_bf8(u.ui32, __amd_scale_to_float(scale), false); +#else + using namespace fcbx; + __amd_fp16x2_storage_t ret; + if (interpret == __AMD_OCP_E4M3) { + ret[0] = to_float<__amd_fp16_storage_t, Encoding::E4M3, true>(val & 0xFF, scale); + ret[1] = to_float<__amd_fp16_storage_t, Encoding::E4M3, true>(val >> 8, scale); + } else { + ret[0] = to_float<__amd_fp16_storage_t, Encoding::E5M2, true>(val & 0xFF, scale); + ret[1] = to_float<__amd_fp16_storage_t, Encoding::E5M2, true>(val >> 8, scale); + } + return ret; +#endif +} + +/** + * @brief Convert fp8 packed 8 to fp16 packed 8 + * + * @param val fp8x8 value + * @param interpret interpretation of fp8 + * @param scale + * @return __amd_fp16x8_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp16x8_storage_t __amd_cvt_fp8x8_to_fp16x8_scale( + const __amd_fp8x8_storage_t val, const __amd_fp8_interpretation_t interpret, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(__amd_fp16x8_storage_t) == sizeof(__amd_fp16x2_storage_t[4])); + static_assert(sizeof(__amd_fp8x8_storage_t) == sizeof(__amd_fp8x2_storage_t[4])); + union { + __amd_fp16x8_storage_t fp16x8; + __amd_fp16x2_storage_t fp16x2[4]; + } ret; + union { + __amd_fp8x8_storage_t fp8x8; + __amd_fp8x2_storage_t fp8x2[4]; + } input{val}; + union { + __amd_fp8x2_storage_t fp8x2[2]; + unsigned int ui32; + } u; + + u.fp8x2[0] = input.fp8x2[0]; + auto adjusted_scale = __amd_scale_to_float(scale); + if (interpret == __AMD_OCP_E4M3) { + ret.fp16x2[0] = __builtin_amdgcn_cvt_scalef32_pk_f16_fp8(u.ui32, adjusted_scale, false); + u.fp8x2[0] = input.fp8x2[1]; + ret.fp16x2[1] = __builtin_amdgcn_cvt_scalef32_pk_f16_fp8(u.ui32, adjusted_scale, false); + u.fp8x2[0] = input.fp8x2[2]; + ret.fp16x2[2] = __builtin_amdgcn_cvt_scalef32_pk_f16_fp8(u.ui32, adjusted_scale, false); + u.fp8x2[0] = input.fp8x2[3]; + ret.fp16x2[3] = __builtin_amdgcn_cvt_scalef32_pk_f16_fp8(u.ui32, adjusted_scale, false); + } else { + ret.fp16x2[0] = __builtin_amdgcn_cvt_scalef32_pk_f16_bf8(u.ui32, adjusted_scale, false); + u.fp8x2[0] = input.fp8x2[1]; + ret.fp16x2[1] = __builtin_amdgcn_cvt_scalef32_pk_f16_bf8(u.ui32, adjusted_scale, false); + u.fp8x2[0] = input.fp8x2[2]; + ret.fp16x2[2] = __builtin_amdgcn_cvt_scalef32_pk_f16_bf8(u.ui32, adjusted_scale, false); + u.fp8x2[0] = input.fp8x2[3]; + ret.fp16x2[3] = __builtin_amdgcn_cvt_scalef32_pk_f16_bf8(u.ui32, adjusted_scale, false); + } + return ret.fp16x8; +#else + using namespace fcbx; + __amd_fp16x8_storage_t ret; + if (interpret == __AMD_OCP_E4M3) { + ret[0] = to_float<__amd_fp16_storage_t, Encoding::E4M3, true>(val[0], scale); + ret[1] = to_float<__amd_fp16_storage_t, Encoding::E4M3, true>(val[1], scale); + ret[2] = to_float<__amd_fp16_storage_t, Encoding::E4M3, true>(val[2], scale); + ret[3] = to_float<__amd_fp16_storage_t, Encoding::E4M3, true>(val[3], scale); + ret[4] = to_float<__amd_fp16_storage_t, Encoding::E4M3, true>(val[4], scale); + ret[5] = to_float<__amd_fp16_storage_t, Encoding::E4M3, true>(val[5], scale); + ret[6] = to_float<__amd_fp16_storage_t, Encoding::E4M3, true>(val[6], scale); + ret[7] = to_float<__amd_fp16_storage_t, Encoding::E4M3, true>(val[7], scale); + } else { + ret[0] = to_float<__amd_fp16_storage_t, Encoding::E5M2, true>(val[0], scale); + ret[1] = to_float<__amd_fp16_storage_t, Encoding::E5M2, true>(val[1], scale); + ret[2] = to_float<__amd_fp16_storage_t, Encoding::E5M2, true>(val[2], scale); + ret[3] = to_float<__amd_fp16_storage_t, Encoding::E5M2, true>(val[3], scale); + ret[4] = to_float<__amd_fp16_storage_t, Encoding::E5M2, true>(val[4], scale); + ret[5] = to_float<__amd_fp16_storage_t, Encoding::E5M2, true>(val[5], scale); + ret[6] = to_float<__amd_fp16_storage_t, Encoding::E5M2, true>(val[6], scale); + ret[7] = to_float<__amd_fp16_storage_t, Encoding::E5M2, true>(val[7], scale); + } + return ret; +#endif +} + +/** + * @brief Convert fp8x2 to bf16x2 with scale. This is direct mapping of gfx950 builtin. + * + * @param in fp8x2 input + * @param interpret + * @param scale + * @return __amd_bf16x2_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_bf16x2_storage_t __amd_cvt_fp8x2_to_bf16x2_scale( + const __amd_fp8x2_storage_t in, const __amd_fp8_interpretation_t interpret, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(unsigned int) == sizeof(__amd_fp8x2_storage_t[2])); + union { + __amd_fp8x2_storage_t fp8x2[2]; + unsigned int ui32; + } u; + u.fp8x2[0] = in; + return interpret == __AMD_OCP_E4M3 + ? __builtin_amdgcn_cvt_scalef32_pk_bf16_fp8(u.ui32, __amd_scale_to_float(scale), false) + : __builtin_amdgcn_cvt_scalef32_pk_bf16_bf8(u.ui32, __amd_scale_to_float(scale), false); +#else + using namespace fcbx; + __amd_bf16x2_storage_t ret; + if (interpret == __AMD_OCP_E4M3) { + ret[0] = to_float<__amd_bf16_storage_t, Encoding::E4M3, true>(in & 0xFF, scale); + ret[1] = to_float<__amd_bf16_storage_t, Encoding::E4M3, true>(in >> 8, scale); + } else { + ret[0] = to_float<__amd_bf16_storage_t, Encoding::E5M2, true>(in & 0xFF, scale); + ret[1] = to_float<__amd_bf16_storage_t, Encoding::E5M2, true>(in >> 8, scale); + } + return ret; +#endif +} + +/** + * @brief Convert fp8 packed 8 to bf16 packed 8. + * + * @param val fp8x8 value + * @param interpret interpretation of fp8 + * @param scale + * @return __amd_bf16x8_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_bf16x8_storage_t __amd_cvt_fp8x8_to_bf16x8_scale( + const __amd_fp8x8_storage_t val, const __amd_fp8_interpretation_t interpret, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(__amd_bf16x8_storage_t) == sizeof(__amd_bf16x2_storage_t[4])); + static_assert(sizeof(__amd_fp8x8_storage_t) == sizeof(__amd_fp8x2_storage_t[4])); + union { + __amd_bf16x8_storage_t bf16x8; + __amd_bf16x2_storage_t bf16x2[4]; + } ret; + union { + __amd_fp8x8_storage_t fp8x8; + __amd_fp8x2_storage_t fp8x2[4]; + } input{val}; + union { + __amd_fp8x2_storage_t fp8x2[2]; + unsigned int ui32; + } u; + + auto adjusted_scale = __amd_scale_to_float(scale); + if (interpret == __AMD_OCP_E4M3) { + u.fp8x2[0] = input.fp8x2[0]; + ret.bf16x2[0] = __builtin_amdgcn_cvt_scalef32_pk_bf16_fp8(u.ui32, adjusted_scale, false); + u.fp8x2[0] = input.fp8x2[1]; + ret.bf16x2[1] = __builtin_amdgcn_cvt_scalef32_pk_bf16_fp8(u.ui32, adjusted_scale, false); + u.fp8x2[0] = input.fp8x2[2]; + ret.bf16x2[2] = __builtin_amdgcn_cvt_scalef32_pk_bf16_fp8(u.ui32, adjusted_scale, false); + u.fp8x2[0] = input.fp8x2[3]; + ret.bf16x2[3] = __builtin_amdgcn_cvt_scalef32_pk_bf16_fp8(u.ui32, adjusted_scale, false); + } else { + u.fp8x2[0] = input.fp8x2[0]; + ret.bf16x2[0] = __builtin_amdgcn_cvt_scalef32_pk_bf16_bf8(u.ui32, adjusted_scale, false); + u.fp8x2[0] = input.fp8x2[1]; + ret.bf16x2[1] = __builtin_amdgcn_cvt_scalef32_pk_bf16_bf8(u.ui32, adjusted_scale, false); + u.fp8x2[0] = input.fp8x2[2]; + ret.bf16x2[2] = __builtin_amdgcn_cvt_scalef32_pk_bf16_bf8(u.ui32, adjusted_scale, false); + u.fp8x2[0] = input.fp8x2[3]; + ret.bf16x2[3] = __builtin_amdgcn_cvt_scalef32_pk_bf16_bf8(u.ui32, adjusted_scale, false); + } + return ret.bf16x8; +#else + using namespace fcbx; + __amd_bf16x8_storage_t ret; + if (interpret == __AMD_OCP_E4M3) { + ret[0] = to_float<__amd_bf16_storage_t, Encoding::E4M3, true>(val[0], scale); + ret[1] = to_float<__amd_bf16_storage_t, Encoding::E4M3, true>(val[1], scale); + ret[2] = to_float<__amd_bf16_storage_t, Encoding::E4M3, true>(val[2], scale); + ret[3] = to_float<__amd_bf16_storage_t, Encoding::E4M3, true>(val[3], scale); + ret[4] = to_float<__amd_bf16_storage_t, Encoding::E4M3, true>(val[4], scale); + ret[5] = to_float<__amd_bf16_storage_t, Encoding::E4M3, true>(val[5], scale); + ret[6] = to_float<__amd_bf16_storage_t, Encoding::E4M3, true>(val[6], scale); + ret[7] = to_float<__amd_bf16_storage_t, Encoding::E4M3, true>(val[7], scale); + } else { + ret[0] = to_float<__amd_bf16_storage_t, Encoding::E5M2, true>(val[0], scale); + ret[1] = to_float<__amd_bf16_storage_t, Encoding::E5M2, true>(val[1], scale); + ret[2] = to_float<__amd_bf16_storage_t, Encoding::E5M2, true>(val[2], scale); + ret[3] = to_float<__amd_bf16_storage_t, Encoding::E5M2, true>(val[3], scale); + ret[4] = to_float<__amd_bf16_storage_t, Encoding::E5M2, true>(val[4], scale); + ret[5] = to_float<__amd_bf16_storage_t, Encoding::E5M2, true>(val[5], scale); + ret[6] = to_float<__amd_bf16_storage_t, Encoding::E5M2, true>(val[6], scale); + ret[7] = to_float<__amd_bf16_storage_t, Encoding::E5M2, true>(val[7], scale); + } + return ret; +#endif +} + +/** + * @brief Convert packed fp6x32 to fp16x32 with scale. + * + * @param in fp6x32 value + * @param interpret + * @param scale + * @return __amd_fp16x32_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp16x32_storage_t __amd_cvt_fp6x32_to_fp16x32_scale( + const __amd_fp6x32_storage_t in, const __amd_fp6_interpretation_t interpret, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + // gfx950 expects scale to be in float + return interpret == __AMD_OCP_E2M3 + ? __builtin_amdgcn_cvt_scalef32_pk32_f16_fp6(in, __amd_scale_to_float(scale)) + : __builtin_amdgcn_cvt_scalef32_pk32_f16_bf6(in, __amd_scale_to_float(scale)); +#else + using namespace fcbx; + if (interpret == __AMD_OCP_E2M3) { + return fp6_cvt_packedx32<__amd_fp6x32_storage_t, __amd_fp16x32_storage_t, __amd_fp16_storage_t, + Encoding::E2M3, Encoding::E5M10>(in, scale); + } else { + return fp6_cvt_packedx32<__amd_fp6x32_storage_t, __amd_fp16x32_storage_t, __amd_fp16_storage_t, + Encoding::E3M2, Encoding::E5M10>(in, scale); + } +#endif +} + +/** + * @brief Convert packed fp6x32 to bf16x32 with scale. + * + * @param in fp6x32 value + * @param interpret + * @param scale + * @return __amd_bf16x32_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_bf16x32_storage_t __amd_cvt_fp6x32_to_bf16x32_scale( + const __amd_fp6x32_storage_t in, const __amd_fp6_interpretation_t interpret, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + return interpret == __AMD_OCP_E2M3 + ? __builtin_amdgcn_cvt_scalef32_pk32_bf16_fp6(in, __amd_scale_to_float(scale)) + : __builtin_amdgcn_cvt_scalef32_pk32_bf16_bf6(in, __amd_scale_to_float(scale)); +#else + using namespace fcbx; + if (interpret == __AMD_OCP_E2M3) { + return fp6_cvt_packedx32<__amd_fp6x32_storage_t, __amd_bf16x32_storage_t, __amd_bf16_storage_t, + Encoding::E2M3, Encoding::E8M7>(in, scale); + } else { + return fp6_cvt_packedx32<__amd_fp6x32_storage_t, __amd_bf16x32_storage_t, __amd_bf16_storage_t, + Encoding::E3M2, Encoding::E8M7>(in, scale); + } +#endif +} + +__OCP_FP_HOST_DEVICE_STATIC__ __amd_floatx32_storage_t __amd_cvt_fp6x32_to_floatx32_scale( + const __amd_fp6x32_storage_t val, const __amd_fp6_interpretation_t interpret, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + return interpret == __AMD_OCP_E2M3 + ? __builtin_amdgcn_cvt_scalef32_pk32_f32_fp6(val, __amd_scale_to_float(scale)) + : __builtin_amdgcn_cvt_scalef32_pk32_f32_bf6(val, __amd_scale_to_float(scale)); +#else + using namespace fcbx; + return interpret == __AMD_OCP_E2M3 + ? fp6_cvt_packedx32<__amd_fp6x32_storage_t, __amd_floatx32_storage_t, float, Encoding::E2M3, + Encoding::IEEE754>(val, scale) + : fp6_cvt_packedx32<__amd_fp6x32_storage_t, __amd_floatx32_storage_t, float, Encoding::E3M2, + Encoding::IEEE754>(val, scale); +#endif +} + +/** + * @brief Convert packed 2 of fp4 to fp16. + * + * @param in packed fp4x2 + * @param scale + * @return __amd_fp16x2_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp16x2_storage_t __amd_cvt_fp4x2_to_fp16x2_scale( + const __amd_fp4x2_storage_t in, const __amd_fp4_interpretation_t, const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + return __builtin_amdgcn_cvt_scalef32_pk_f16_fp4(in, __amd_scale_to_float(scale), 0); +#else + using namespace fcbx; + __amd_fp16x2_storage_t ret{to_float<__amd_fp16_storage_t, Encoding::E2M1, true>(in & 0xFu, scale), + to_float<__amd_fp16_storage_t, Encoding::E2M1, true>(in >> 4, scale)}; + return ret; +#endif +} + +/** + * @brief convert packed fp4x8 to fp16 x16. + * + * @param in + * @param scale + * @return __amd_fp16x8_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp16x8_storage_t __amd_cvt_fp4x8_to_fp16x8_scale( + const __amd_fp4x8_storage_t in, const __amd_fp4_interpretation_t, const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(__amd_fp4x2_storage_t[4]) == sizeof(unsigned int)); + static_assert(sizeof(__amd_fp16x2_storage_t[4]) == sizeof(__amd_fp16x8_storage_t)); + union { + __amd_fp4x8_storage_t fp4x8; + __amd_fp4x2_storage_t fp4x2[4]; + } u{in}; + union { + __amd_fp16x8_storage_t fp16x8; + __amd_fp16x2_storage_t fp16x2[4]; + } ret; + ret.fp16x2[0] = + __builtin_amdgcn_cvt_scalef32_pk_f16_fp4(u.fp4x2[0], __amd_scale_to_float(scale), 0); + ret.fp16x2[1] = + __builtin_amdgcn_cvt_scalef32_pk_f16_fp4(u.fp4x2[1], __amd_scale_to_float(scale), 0); + ret.fp16x2[2] = + __builtin_amdgcn_cvt_scalef32_pk_f16_fp4(u.fp4x2[2], __amd_scale_to_float(scale), 0); + ret.fp16x2[3] = + __builtin_amdgcn_cvt_scalef32_pk_f16_fp4(u.fp4x2[3], __amd_scale_to_float(scale), 0); + return ret.fp16x8; +#else + using namespace fcbx; + __amd_fp16x8_storage_t ret; + ret[0] = to_float<__amd_fp16_storage_t, Encoding::E2M1, true>(in & 0xFu, scale); + ret[1] = to_float<__amd_fp16_storage_t, Encoding::E2M1, true>((in >> 4) & 0xFu, scale); + ret[2] = to_float<__amd_fp16_storage_t, Encoding::E2M1, true>((in >> 8) & 0xFu, scale); + ret[3] = to_float<__amd_fp16_storage_t, Encoding::E2M1, true>((in >> 12) & 0xFu, scale); + ret[4] = to_float<__amd_fp16_storage_t, Encoding::E2M1, true>((in >> 16) & 0xFu, scale); + ret[5] = to_float<__amd_fp16_storage_t, Encoding::E2M1, true>((in >> 20) & 0xFu, scale); + ret[6] = to_float<__amd_fp16_storage_t, Encoding::E2M1, true>((in >> 24) & 0xFu, scale); + ret[7] = to_float<__amd_fp16_storage_t, Encoding::E2M1, true>((in >> 28) & 0xFu, scale); + return ret; +#endif +} + +/** + * @brief Convert packed fp4x2 to bf16x2. + * + * @param in + * @param scale + * @return __OCP_FP_DEVICE_STATIC__ + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_bf16x2_storage_t __amd_cvt_fp4x2_to_bf16x2_scale( + const __amd_fp4x2_storage_t in, const __amd_fp4_interpretation_t, const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + return __builtin_amdgcn_cvt_scalef32_pk_bf16_fp4(in, __amd_scale_to_float(scale), 0); +#else + using namespace fcbx; + __amd_bf16x2_storage_t ret{to_float<__amd_bf16_storage_t, Encoding::E2M1, true>(in & 0xFu, scale), + to_float<__amd_bf16_storage_t, Encoding::E2M1, true>(in >> 4, scale)}; + return ret; +#endif +} + +/** + * @brief Convert packed fp4x8 to bf16x8. + * + * @param in + * @param scale + * @return __amd_bf16x8_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_bf16x8_storage_t __amd_cvt_fp4x8_to_bf16x8_scale( + const __amd_fp4x8_storage_t in, const __amd_fp4_interpretation_t, const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(__amd_fp4x2_storage_t[4]) == sizeof(unsigned int)); + static_assert(sizeof(__amd_bf16x2_storage_t[4]) == sizeof(__amd_bf16x8_storage_t)); + union { + __amd_fp4x8_storage_t fp4x8; + __amd_fp4x2_storage_t fp4x2[4]; + } u{in}; + union { + __amd_bf16x8_storage_t bf16x8; + __amd_bf16x2_storage_t bf16x2[4]; + } ret; + // Unrolled + ret.bf16x2[0] = + __builtin_amdgcn_cvt_scalef32_pk_bf16_fp4(u.fp4x2[0], __amd_scale_to_float(scale), 0); + ret.bf16x2[1] = + __builtin_amdgcn_cvt_scalef32_pk_bf16_fp4(u.fp4x2[1], __amd_scale_to_float(scale), 0); + ret.bf16x2[2] = + __builtin_amdgcn_cvt_scalef32_pk_bf16_fp4(u.fp4x2[2], __amd_scale_to_float(scale), 0); + ret.bf16x2[3] = + __builtin_amdgcn_cvt_scalef32_pk_bf16_fp4(u.fp4x2[3], __amd_scale_to_float(scale), 0); + return ret.bf16x8; +#else + using namespace fcbx; + __amd_bf16x8_storage_t ret; + ret[0] = to_float<__amd_bf16_storage_t, Encoding::E2M1, true>(in & 0xFu, scale); + ret[1] = to_float<__amd_bf16_storage_t, Encoding::E2M1, true>((in >> 4) & 0xFu, scale); + ret[2] = to_float<__amd_bf16_storage_t, Encoding::E2M1, true>((in >> 8) & 0xFu, scale); + ret[3] = to_float<__amd_bf16_storage_t, Encoding::E2M1, true>((in >> 12) & 0xFu, scale); + ret[4] = to_float<__amd_bf16_storage_t, Encoding::E2M1, true>((in >> 16) & 0xFu, scale); + ret[5] = to_float<__amd_bf16_storage_t, Encoding::E2M1, true>((in >> 20) & 0xFu, scale); + ret[6] = to_float<__amd_bf16_storage_t, Encoding::E2M1, true>((in >> 24) & 0xFu, scale); + ret[7] = to_float<__amd_bf16_storage_t, Encoding::E2M1, true>((in >> 28) & 0xFu, scale); + return ret; +#endif +} + +/** + * @brief Convert packed fp4x8 to floatx8 with scale. + * + * @param val + * @param scale + * @return __amd_floatx8_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_floatx8_storage_t __amd_cvt_fp4x8_to_floatx8_scale( + const __amd_fp4x8_storage_t val, const __amd_fp4_interpretation_t, const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(__amd_fp4x2_storage_t[4]) == sizeof(__amd_fp4x8_storage_t)); + union { + __amd_fp4x8_storage_t fp4x8; + __amd_fp4x2_storage_t fp8x2[4]; + } u{val}; + __amd_floatx2_storage_t op; + __amd_floatx8_storage_t ret; + op = __builtin_amdgcn_cvt_scalef32_pk_f32_fp4(u.fp8x2[0], __amd_scale_to_float(scale), 0); + ret[0] = op[0]; + ret[1] = op[1]; + op = __builtin_amdgcn_cvt_scalef32_pk_f32_fp4(u.fp8x2[1], __amd_scale_to_float(scale), 0); + ret[2] = op[0]; + ret[3] = op[1]; + op = __builtin_amdgcn_cvt_scalef32_pk_f32_fp4(u.fp8x2[2], __amd_scale_to_float(scale), 0); + ret[4] = op[0]; + ret[5] = op[1]; + op = __builtin_amdgcn_cvt_scalef32_pk_f32_fp4(u.fp8x2[3], __amd_scale_to_float(scale), 0); + ret[6] = op[0]; + ret[7] = op[1]; + return ret; +#else + using namespace fcbx; + __amd_floatx8_storage_t ret; + ret[0] = to_float(val & 0xFu, scale); + ret[1] = to_float((val >> 4) & 0xFu, scale); + ret[2] = to_float((val >> 8) & 0xFu, scale); + ret[3] = to_float((val >> 12) & 0xFu, scale); + ret[4] = to_float((val >> 16) & 0xFu, scale); + ret[5] = to_float((val >> 20) & 0xFu, scale); + ret[6] = to_float((val >> 24) & 0xFu, scale); + ret[7] = to_float((val >> 28) & 0xFu, scale); + return ret; +#endif +} + +/** + * @brief Convert packed floatx8 to fp4x8. + * + * @param in + * @param interpret + * @param scale + * @return __amd_fp4x8_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp4x8_storage_t __amd_cvt_floatx8_to_fp4x8_scale( + const __amd_floatx8_storage_t in, const __amd_fp4_interpretation_t, const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(unsigned int) == sizeof(__amd_fp4x2_storage_t[4])); + static_assert(sizeof(__amd_fp4x8_storage_t) == sizeof(__amd_fp4x2_storage_t[4])); + union hold_u { + unsigned int ui32; + __amd_fp4x2_storage_t fp4x2[4]; + __amd_fp4x8_storage_t fp4x8; + } ret{0}, tmp{0}; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_pk_fp4_f32(tmp.ui32, in[0], in[1], + __amd_scale_to_float(scale), 0); + ret.fp4x2[0] = tmp.fp4x2[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_pk_fp4_f32(tmp.ui32, in[2], in[3], + __amd_scale_to_float(scale), 0); + ret.fp4x2[1] = tmp.fp4x2[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_pk_fp4_f32(tmp.ui32, in[4], in[5], + __amd_scale_to_float(scale), 0); + ret.fp4x2[2] = tmp.fp4x2[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_pk_fp4_f32(tmp.ui32, in[6], in[7], + __amd_scale_to_float(scale), 0); + ret.fp4x2[3] = tmp.fp4x2[0]; + return ret.fp4x8; +#else + __amd_fp4x8_storage_t ret = 0; + using namespace fcbx; + auto tmp = from_float(in[7], scale); + ret |= tmp; + tmp = from_float(in[6], scale); + ret <<= 4; + ret |= tmp; + tmp = from_float(in[5], scale); + ret <<= 4; + ret |= tmp; + tmp = from_float(in[4], scale); + ret <<= 4; + ret |= tmp; + tmp = from_float(in[3], scale); + ret <<= 4; + ret |= tmp; + tmp = from_float(in[2], scale); + ret <<= 4; + ret |= tmp; + tmp = from_float(in[1], scale); + ret <<= 4; + ret |= tmp; + tmp = from_float(in[0], scale); + ret <<= 4; + ret |= tmp; + return ret; +#endif +} + +/** + * @brief Convert packed fp16x2 to fp8x2 with scale. + * + * @param in + * @param interpret + * @param scale + * @return __amd_fp8x2_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp8x2_storage_t __amd_cvt_fp16x2_to_fp8x2_scale( + const __amd_fp16x2_storage_t in, const __amd_fp8_interpretation_t interpret, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(__amd_shortx2_storage_t) == sizeof(__amd_fp8x2_storage_t[2])); + static_assert(sizeof(uint32_t) == sizeof(__amd_fp8x2_storage_t[2])); + union { + __amd_shortx2_storage_t shortx2; + __amd_fp8x2_storage_t fp8x2[2]; + } u{0}; + u.shortx2 = interpret == __AMD_OCP_E4M3 + ? __builtin_amdgcn_cvt_scalef32_pk_fp8_f16(u.shortx2, in, __amd_scale_to_float(scale), false) + : __builtin_amdgcn_cvt_scalef32_pk_bf8_f16(u.shortx2, in, __amd_scale_to_float(scale), false); + return u.fp8x2[0]; +#else + static_assert(sizeof(__amd_fp8x2_storage_t[2]) == sizeof(uint32_t)); + using namespace fcbx; + union { + uint32_t ui32; + __amd_fp8x2_storage_t fp8x2[2]; + } u{0}; + if (interpret == __AMD_OCP_E4M3) { + u.ui32 = from_float<__amd_fp16_storage_t, Encoding::E4M3, true>(in[1], scale); + u.ui32 <<= 8; + u.ui32 |= from_float<__amd_fp16_storage_t, Encoding::E4M3, true>(in[0], scale); + } else { + u.ui32 = from_float<__amd_fp16_storage_t, Encoding::E5M2, true>(in[1], scale); + u.ui32 <<= 8; + u.ui32 |= from_float<__amd_fp16_storage_t, Encoding::E5M2, true>(in[0], scale); + } + return u.fp8x2[0]; +#endif +} + +/** + * @brief Convert packed bf16x2 to fp8x2 with scale. + * + * @param in + * @param interpret + * @param scale + * @return __amd_fp8x2_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp8x2_storage_t __amd_cvt_bf16x2_to_fp8x2_scale( + const __amd_bf16x2_storage_t in, const __amd_fp8_interpretation_t interpret, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(__amd_shortx2_storage_t) == sizeof(__amd_fp8x2_storage_t[2])); + union { + __amd_shortx2_storage_t shortx2; + __amd_fp8x2_storage_t fp8x2[2]; + } u{0}; + u.shortx2 = interpret == __AMD_OCP_E4M3 + ? __builtin_amdgcn_cvt_scalef32_pk_fp8_bf16(u.shortx2, in, __amd_scale_to_float(scale), false) + : __builtin_amdgcn_cvt_scalef32_pk_bf8_bf16(u.shortx2, in, __amd_scale_to_float(scale), + false); + return u.fp8x2[0]; +#else + using namespace fcbx; + union { + uint32_t ui32; + __amd_fp8x2_storage_t fp8x2[2]; + } u{0}; + if (interpret == __AMD_OCP_E4M3) { + u.ui32 = from_float<__amd_bf16_storage_t, Encoding::E4M3, true>(in[1], scale); + u.ui32 <<= 8; + u.ui32 |= from_float<__amd_bf16_storage_t, Encoding::E4M3, true>(in[0], scale); + } else { + u.ui32 = from_float<__amd_bf16_storage_t, Encoding::E5M2, true>(in[1], scale); + u.ui32 <<= 8; + u.ui32 |= from_float<__amd_bf16_storage_t, Encoding::E5M2, true>(in[0], scale); + } + return u.fp8x2[0]; +#endif +} + +/** + * @brief Convert bf16 pack 8 to fp8 packed 8 + * + * @param val bf16x8 value + * @param interpret interpretation of fp8 + * @param scale + * @return __amd_fp8x8_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp8x8_storage_t __amd_cvt_bf16x8_to_fp8x8_scale( + const __amd_bf16x8_storage_t val, const __amd_fp8_interpretation_t interpret, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(__amd_fp8x2_storage_t[4]) == sizeof(__amd_fp8x8_storage_t)); + static_assert(sizeof(__amd_fp8x2_storage_t[2]) == sizeof(unsigned int)); + union { + __amd_shortx2_storage_t shortx2; + __amd_fp8x2_storage_t fp8x2[2]; + unsigned int ui32; + } u{0}; + union { + __amd_fp8x2_storage_t fp8x2[4]; + __amd_fp8x8_storage_t fp8x8; + } result; + __amd_shortx2_storage_t t_shortx2{0, 0}; + if (interpret == __AMD_OCP_E4M3) { + u.shortx2 = __builtin_amdgcn_cvt_scalef32_pk_fp8_bf16( + t_shortx2, __amd_bf16x2_storage_t{val[0], val[1]}, __amd_scale_to_float(scale), false); + result.fp8x2[0] = u.fp8x2[0]; + u.shortx2 = __builtin_amdgcn_cvt_scalef32_pk_fp8_bf16( + t_shortx2, __amd_bf16x2_storage_t{val[2], val[3]}, __amd_scale_to_float(scale), false); + result.fp8x2[1] = u.fp8x2[0]; + u.shortx2 = __builtin_amdgcn_cvt_scalef32_pk_fp8_bf16( + t_shortx2, __amd_bf16x2_storage_t{val[4], val[5]}, __amd_scale_to_float(scale), false); + result.fp8x2[2] = u.fp8x2[0]; + u.shortx2 = __builtin_amdgcn_cvt_scalef32_pk_fp8_bf16( + t_shortx2, __amd_bf16x2_storage_t{val[6], val[7]}, __amd_scale_to_float(scale), false); + result.fp8x2[3] = u.fp8x2[0]; + } else { + u.shortx2 = __builtin_amdgcn_cvt_scalef32_pk_bf8_bf16( + t_shortx2, __amd_bf16x2_storage_t{val[0], val[1]}, __amd_scale_to_float(scale), false); + result.fp8x2[0] = u.fp8x2[0]; + u.shortx2 = __builtin_amdgcn_cvt_scalef32_pk_bf8_bf16( + t_shortx2, __amd_bf16x2_storage_t{val[2], val[3]}, __amd_scale_to_float(scale), false); + result.fp8x2[1] = u.fp8x2[0]; + u.shortx2 = __builtin_amdgcn_cvt_scalef32_pk_bf8_bf16( + t_shortx2, __amd_bf16x2_storage_t{val[4], val[5]}, __amd_scale_to_float(scale), false); + result.fp8x2[2] = u.fp8x2[0]; + u.shortx2 = __builtin_amdgcn_cvt_scalef32_pk_bf8_bf16( + t_shortx2, __amd_bf16x2_storage_t{val[6], val[7]}, __amd_scale_to_float(scale), false); + result.fp8x2[3] = u.fp8x2[0]; + } + return result.fp8x8; +#else + using namespace fcbx; + __amd_fp8x8_storage_t ret; + if (interpret == __AMD_OCP_E4M3) { + ret[0] = from_float<__amd_bf16_storage_t, Encoding::E4M3, true>(val[0], scale); + ret[1] = from_float<__amd_bf16_storage_t, Encoding::E4M3, true>(val[1], scale); + ret[2] = from_float<__amd_bf16_storage_t, Encoding::E4M3, true>(val[2], scale); + ret[3] = from_float<__amd_bf16_storage_t, Encoding::E4M3, true>(val[3], scale); + ret[4] = from_float<__amd_bf16_storage_t, Encoding::E4M3, true>(val[4], scale); + ret[5] = from_float<__amd_bf16_storage_t, Encoding::E4M3, true>(val[5], scale); + ret[6] = from_float<__amd_bf16_storage_t, Encoding::E4M3, true>(val[6], scale); + ret[7] = from_float<__amd_bf16_storage_t, Encoding::E4M3, true>(val[7], scale); + } else { + ret[0] = from_float<__amd_bf16_storage_t, Encoding::E5M2, true>(val[0], scale); + ret[1] = from_float<__amd_bf16_storage_t, Encoding::E5M2, true>(val[1], scale); + ret[2] = from_float<__amd_bf16_storage_t, Encoding::E5M2, true>(val[2], scale); + ret[3] = from_float<__amd_bf16_storage_t, Encoding::E5M2, true>(val[3], scale); + ret[4] = from_float<__amd_bf16_storage_t, Encoding::E5M2, true>(val[4], scale); + ret[5] = from_float<__amd_bf16_storage_t, Encoding::E5M2, true>(val[5], scale); + ret[6] = from_float<__amd_bf16_storage_t, Encoding::E5M2, true>(val[6], scale); + ret[7] = from_float<__amd_bf16_storage_t, Encoding::E5M2, true>(val[7], scale); + } + return ret; +#endif +} + +/** + * @brief Convert fp8 packed 8 to float packed 8. + * + * @param val fp8x8 value + * @param interpret interpretation of fp8 + * @param scale + * @return __amd_floatx8_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_floatx8_storage_t __amd_cvt_fp8x8_to_floatx8_scale( + const __amd_fp8x8_storage_t val, const __amd_fp8_interpretation_t interpret, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(__amd_fp8x8_storage_t) == sizeof(__amd_fp8x2_storage_t[4])); + union { + __amd_fp8x8_storage_t fp8x8; + __amd_fp8x2_storage_t fp8x2[4]; + } u{val}; + __amd_floatx8_storage_t ret; + __amd_floatx2_storage_t out; + if (interpret == __AMD_OCP_E4M3) { + out = __builtin_amdgcn_cvt_scalef32_pk_f32_fp8(u.fp8x2[0], __amd_scale_to_float(scale), false); + ret[0] = out[0]; + ret[1] = out[1]; + out = __builtin_amdgcn_cvt_scalef32_pk_f32_fp8(u.fp8x2[1], __amd_scale_to_float(scale), false); + ret[2] = out[0]; + ret[3] = out[1]; + out = __builtin_amdgcn_cvt_scalef32_pk_f32_fp8(u.fp8x2[2], __amd_scale_to_float(scale), false); + ret[4] = out[0]; + ret[5] = out[1]; + out = __builtin_amdgcn_cvt_scalef32_pk_f32_fp8(u.fp8x2[3], __amd_scale_to_float(scale), false); + ret[6] = out[0]; + ret[7] = out[1]; + } else { + out = __builtin_amdgcn_cvt_scalef32_pk_f32_bf8(u.fp8x2[0], __amd_scale_to_float(scale), false); + ret[0] = out[0]; + ret[1] = out[1]; + out = __builtin_amdgcn_cvt_scalef32_pk_f32_bf8(u.fp8x2[1], __amd_scale_to_float(scale), false); + ret[2] = out[0]; + ret[3] = out[1]; + out = __builtin_amdgcn_cvt_scalef32_pk_f32_bf8(u.fp8x2[2], __amd_scale_to_float(scale), false); + ret[4] = out[0]; + ret[5] = out[1]; + out = __builtin_amdgcn_cvt_scalef32_pk_f32_bf8(u.fp8x2[3], __amd_scale_to_float(scale), false); + ret[6] = out[0]; + ret[7] = out[1]; + } + return ret; +#else + using namespace fcbx; + __amd_floatx8_storage_t ret; + if (interpret == __AMD_OCP_E4M3) { + ret[0] = to_float(val[0], scale); + ret[1] = to_float(val[1], scale); + ret[2] = to_float(val[2], scale); + ret[3] = to_float(val[3], scale); + ret[4] = to_float(val[4], scale); + ret[5] = to_float(val[5], scale); + ret[6] = to_float(val[6], scale); + ret[7] = to_float(val[7], scale); + } else { + ret[0] = to_float(val[0], scale); + ret[1] = to_float(val[1], scale); + ret[2] = to_float(val[2], scale); + ret[3] = to_float(val[3], scale); + ret[4] = to_float(val[4], scale); + ret[5] = to_float(val[5], scale); + ret[6] = to_float(val[6], scale); + ret[7] = to_float(val[7], scale); + } + return ret; +#endif +} + +/** + * @brief Convert fp8 to fp16 with scale. + * + * @param val + * @param interpret + * @param scale + * @return __amd_fp16_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp16_storage_t +__amd_cvt_fp8_to_fp16_scale(const __amd_fp8_storage_t val, + const __amd_fp8_interpretation_t interpret, const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + __amd_fp16x2_storage_t ret; + ret = interpret == __AMD_OCP_E4M3 + ? __builtin_amdgcn_cvt_scalef32_f16_fp8(ret, val, __amd_scale_to_float(scale), 0, false) + : __builtin_amdgcn_cvt_scalef32_f16_bf8(ret, val, __amd_scale_to_float(scale), 0, false); + return ret[0]; +#else + using namespace fcbx; + if (interpret == __AMD_OCP_E4M3) { + return to_float<__amd_fp16_storage_t, Encoding::E4M3, true>(val, scale); + } else { + return to_float<__amd_fp16_storage_t, Encoding::E5M2, true>(val, scale); + } +#endif +} + +/** + * @brief Convert fp8 to bf16 with scale. + * + * @param val + * @param interpret + * @param scale + * @return __amd_bf16_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_bf16_storage_t +__amd_cvt_fp8_to_bf16_scale(const __amd_fp8_storage_t val, + const __amd_fp8_interpretation_t interpret, const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(unsigned int) == sizeof(__amd_fp8x2_storage_t[2])); + static_assert(sizeof(__amd_fp8_storage_t[4]) == sizeof(__amd_fp8x2_storage_t[2])); + union { + __amd_fp8_storage_t fp8[4]; + __amd_fp8x2_storage_t fp8x2[2]; + unsigned int ui32; + } u{0}; + u.fp8[0] = val; + auto ret = interpret == __AMD_OCP_E4M3 + ? __builtin_amdgcn_cvt_scalef32_pk_bf16_fp8(u.ui32, __amd_scale_to_float(scale), false) + : __builtin_amdgcn_cvt_scalef32_pk_bf16_bf8(u.ui32, __amd_scale_to_float(scale), false); + return ret[0]; +#else + using namespace fcbx; + if (interpret == __AMD_OCP_E4M3) { + return to_float<__amd_bf16_storage_t, Encoding::E4M3, true>(val, scale); + } else { + return to_float<__amd_bf16_storage_t, Encoding::E5M2, true>(val, scale); + } +#endif +} + +/** + * @brief Convert two packed float16x16 to fp6x32. + * + * @param in1 + * @param in2 + * @param interpret + * @param scale + * @return __amd_fp6x32_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp6x32_storage_t __amd_cvt_floatx16_floatx16_to_fp6x32_scale( + const __amd_floatx16_storage_t in1, const __amd_floatx16_storage_t in2, + const __amd_fp6_interpretation_t interpret, const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + return interpret == __AMD_OCP_E2M3 + ? __builtin_amdgcn_cvt_scalef32_2xpk16_fp6_f32(in1, in2, __amd_scale_to_float(scale)) + : __builtin_amdgcn_cvt_scalef32_2xpk16_bf6_f32(in1, in2, __amd_scale_to_float(scale)); +#else + __amd_floatx32_storage_t tmp; + for (size_t i = 0; i < 16; i++) { + tmp[i] = in1[i]; + } + for (size_t i = 0; i < 16; i++) { + tmp[i + 16] = in2[i]; + } + using namespace fcbx; + return interpret == __AMD_OCP_E2M3 + ? fp6_cvt_packedx32<__amd_floatx32_storage_t, __amd_fp6x32_storage_t, float, + Encoding::IEEE754, Encoding::E2M3>(tmp, scale) + : fp6_cvt_packedx32<__amd_floatx32_storage_t, __amd_fp6x32_storage_t, float, + Encoding::IEEE754, Encoding::E3M2>(tmp, scale); +#endif +} + +/** + * @brief Convert packed floatx32 to fp6x32. + * + * @param in1 + * @param in2 + * @param interpret + * @param scale + * @return __amd_fp6x32_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp6x32_storage_t __amd_cvt_floatx32_to_fp6x32_scale( + const __amd_floatx32_storage_t val, const __amd_fp6_interpretation_t interpret, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + __amd_floatx16_storage_t in1{val[0], val[1], val[2], val[3], val[4], val[5], + val[6], val[7], val[8], val[9], val[10], val[11], + val[12], val[13], val[14], val[15]}, + in2 = {val[16], val[17], val[18], val[19], val[20], val[21], val[22], val[23], + val[24], val[25], val[26], val[27], val[28], val[29], val[30], val[31]}; + return interpret == __AMD_OCP_E2M3 + ? __builtin_amdgcn_cvt_scalef32_2xpk16_fp6_f32(in1, in2, __amd_scale_to_float(scale)) + : __builtin_amdgcn_cvt_scalef32_2xpk16_bf6_f32(in1, in2, __amd_scale_to_float(scale)); +#else + using namespace fcbx; + return interpret == __AMD_OCP_E2M3 + ? fp6_cvt_packedx32<__amd_floatx32_storage_t, __amd_fp6x32_storage_t, float, + Encoding::IEEE754, Encoding::E2M3>(val, scale) + : fp6_cvt_packedx32<__amd_floatx32_storage_t, __amd_fp6x32_storage_t, float, + Encoding::IEEE754, Encoding::E3M2>(val, scale); +#endif +} + +/** + * @brief Convert packed floatx32 to fp6x32 with stochastic rounding and scale. + * + * @param val + * @param interpret + * @param round + * @param scale + * @return __amd_fp6x32_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp6x32_storage_t __amd_cvt_floatx32_to_fp6x32_sr_scale( + const __amd_floatx32_storage_t val, const __amd_fp6_interpretation_t interpret, + const unsigned int round, const __amd_scale_t scale) { +#if __has_builtin(__builtin_amdgcn_cvt_scalef32_sr_pk32_fp6_f32) and \ + __has_builtin(__builtin_amdgcn_cvt_scalef32_sr_pk32_bf6_f32) + return interpret == __AMD_OCP_E2M3 + ? __builtin_amdgcn_cvt_scalef32_sr_pk32_fp6_f32(val, round, __amd_scale_to_float(scale)) + : __builtin_amdgcn_cvt_scalef32_sr_pk32_bf6_f32(val, round, __amd_scale_to_float(scale)); +#else + using namespace fcbx; + return interpret == __AMD_OCP_E2M3 + ? fp6_cvt_packedx32<__amd_floatx32_storage_t, __amd_fp6x32_storage_t, float, + Encoding::IEEE754, Encoding::E2M3, true>(val, scale, round) + : fp6_cvt_packedx32<__amd_floatx32_storage_t, __amd_fp6x32_storage_t, float, + Encoding::IEEE754, Encoding::E3M2, true>(val, scale, round); +#endif +} + +/** + * @brief Convert float to fp16 with stochastic rounding. + * + * @param in input float val + * @param round + * @return __amd_fp16_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp16_storage_t +__amd_cvt_float_to_fp16_sr(const float in, const unsigned int round) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + __amd_fp16x2_storage_t ret; + ret = __builtin_amdgcn_cvt_sr_f16_f32(ret, in, round, 0); + return ret[0]; +#else + __builtin_trap(); +#endif +} + +/** + * @brief Convert two float inputs to fp16x2. + * + * @param in1 input float val + * @param in2 input float val + * @param round + * @return __amd_fp16x2_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp16x2_storage_t +__amd_cvt_float_float_to_fp16x2_sr(const float in1, const float in2, const unsigned int round) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + __amd_fp16x2_storage_t ret1, ret2; + ret1 = __builtin_amdgcn_cvt_sr_f16_f32(ret1, in1, round, 0); + ret2 = __builtin_amdgcn_cvt_sr_f16_f32(ret2, in2, round, 0); + return __amd_fp16x2_storage_t{ret1[0], ret2[0]}; +#else + __builtin_trap(); +#endif +} + +/** + * @brief Convert float to bfloat16 with stochastic rounding. + * + * @param in + * @param round + * @return __amd_bf16_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_bf16_storage_t +__amd_cvt_float_to_bf16_sr(const float in, const unsigned int round) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + __amd_bf16x2_storage_t ret; + ret = __builtin_amdgcn_cvt_sr_bf16_f32(ret, in, round, 0); + return ret[0]; +#else + __builtin_trap(); +#endif +} + +/** + * @brief Convert packed fp16x32 to fp6x32 with stochastic rounding and scale. + * + * @param in + * @param interpret + * @param round + * @param scale + * @return __amd_fp6x32_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp6x32_storage_t __amd_cvt_fp16x32_to_fp6x32_sr_scale( + const __amd_fp16x32_storage_t in, const __amd_fp6_interpretation_t interpret, + const unsigned int round, const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + return interpret == __AMD_OCP_E2M3 + ? __builtin_amdgcn_cvt_scalef32_sr_pk32_fp6_f16(in, round, __amd_scale_to_float(scale)) + : __builtin_amdgcn_cvt_scalef32_sr_pk32_bf6_f16(in, round, __amd_scale_to_float(scale)); +#else + return interpret == __AMD_OCP_E2M3 + ? fcbx::fp6_cvt_packedx32<__amd_fp16x32_storage_t, __amd_fp6x32_storage_t, + __amd_fp16_storage_t, fcbx::Encoding::E5M10, fcbx::Encoding::E2M3, + true>(in, scale, round) + : fcbx::fp6_cvt_packedx32<__amd_fp16x32_storage_t, __amd_fp6x32_storage_t, + __amd_fp16_storage_t, fcbx::Encoding::E5M10, fcbx::Encoding::E3M2, + true>(in, scale, round); +#endif +} + +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp6x32_storage_t __amd_cvt_bf16x32_to_fp6x32_sr_scale( + const __amd_bf16x32_storage_t in, const __amd_fp6_interpretation_t interpret, + const unsigned int round, const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + return interpret == __AMD_OCP_E2M3 + ? __builtin_amdgcn_cvt_scalef32_sr_pk32_fp6_bf16(in, round, __amd_scale_to_float(scale)) + : __builtin_amdgcn_cvt_scalef32_sr_pk32_bf6_bf16(in, round, __amd_scale_to_float(scale)); +#else + return interpret == __AMD_OCP_E2M3 + ? fcbx::fp6_cvt_packedx32<__amd_bf16x32_storage_t, __amd_fp6x32_storage_t, + __amd_bf16_storage_t, fcbx::Encoding::E8M7, fcbx::Encoding::E2M3, + true>(in, scale, round) + : fcbx::fp6_cvt_packedx32<__amd_bf16x32_storage_t, __amd_fp6x32_storage_t, + __amd_bf16_storage_t, fcbx::Encoding::E8M7, fcbx::Encoding::E3M2, + true>(in, scale, round); +#endif +} + +/** + * @brief Convert packed bf16x2 to fp4x2 with scale. + * + * @param val + * @param scale + * @return __amd_fp4x2_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp4x2_storage_t __amd_cvt_bf16x2_to_fp4x2_scale( + const __amd_bf16x2_storage_t val, const __amd_fp4_interpretation_t, const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(__amd_fp4x2_storage_t[4]) == sizeof(unsigned int)); + union { + unsigned int ui32; + __amd_fp4x2_storage_t fp4x2[4]; + } u{0}; + u.ui32 = __builtin_amdgcn_cvt_scalef32_pk_fp4_bf16(u.ui32, val, __amd_scale_to_float(scale), 0); + return u.fp4x2[0]; +#else + using namespace fcbx; + __amd_fp4x2_storage_t ret; + ret = from_float<__amd_bf16_storage_t, Encoding::E2M1, true>(val[1], scale); + ret <<= 4; + ret |= from_float<__amd_bf16_storage_t, Encoding::E2M1, true>(val[0], scale); + return ret; +#endif +} + +/** + * @brief Convert packed bf16x8 to fp4x8. + * + * @param val + * @param scale + * @return __amd_fp4x8_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp4x8_storage_t __amd_cvt_bf16x8_to_fp4x8_scale( + const __amd_bf16x8_storage_t val, const __amd_fp4_interpretation_t, const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(__amd_fp4x2_storage_t[4]) == sizeof(unsigned int)); + static_assert(sizeof(__amd_fp4x2_storage_t[4]) == sizeof(__amd_fp4x8_storage_t)); + union hold_u { + unsigned int ui32; + __amd_fp4x2_storage_t fp4x2[4]; + __amd_fp4x8_storage_t fp4x8; + } ret{0}, tmp{0}; + __amd_bf16x2_storage_t tmp_in{val[0], val[1]}; + tmp.ui32 = + __builtin_amdgcn_cvt_scalef32_pk_fp4_bf16(tmp.ui32, tmp_in, __amd_scale_to_float(scale), 0); + ret.fp4x2[0] = tmp.fp4x2[0]; + tmp_in[0] = val[2]; + tmp_in[1] = val[3]; + tmp.ui32 = + __builtin_amdgcn_cvt_scalef32_pk_fp4_bf16(tmp.ui32, tmp_in, __amd_scale_to_float(scale), 0); + ret.fp4x2[1] = tmp.fp4x2[0]; + tmp_in[0] = val[4]; + tmp_in[1] = val[5]; + tmp.ui32 = + __builtin_amdgcn_cvt_scalef32_pk_fp4_bf16(tmp.ui32, tmp_in, __amd_scale_to_float(scale), 0); + ret.fp4x2[2] = tmp.fp4x2[0]; + tmp_in[0] = val[6]; + tmp_in[1] = val[7]; + tmp.ui32 = + __builtin_amdgcn_cvt_scalef32_pk_fp4_bf16(tmp.ui32, tmp_in, __amd_scale_to_float(scale), 0); + ret.fp4x2[3] = tmp.fp4x2[0]; + return ret.fp4x8; +#else + __amd_fp4x8_storage_t ret = 0; + using namespace fcbx; + auto tmp = from_float<__amd_bf16_storage_t, Encoding::E2M1, true>(val[7], scale); + ret <<= 4; + ret |= tmp; + tmp = from_float<__amd_bf16_storage_t, Encoding::E2M1, true>(val[6], scale); + ret <<= 4; + ret |= tmp; + tmp = from_float<__amd_bf16_storage_t, Encoding::E2M1, true>(val[5], scale); + ret <<= 4; + ret |= tmp; + tmp = from_float<__amd_bf16_storage_t, Encoding::E2M1, true>(val[4], scale); + ret <<= 4; + ret |= tmp; + tmp = from_float<__amd_bf16_storage_t, Encoding::E2M1, true>(val[3], scale); + ret <<= 4; + ret |= tmp; + tmp = from_float<__amd_bf16_storage_t, Encoding::E2M1, true>(val[2], scale); + ret <<= 4; + ret |= tmp; + tmp = from_float<__amd_bf16_storage_t, Encoding::E2M1, true>(val[1], scale); + ret <<= 4; + ret |= tmp; + tmp = from_float<__amd_bf16_storage_t, Encoding::E2M1, true>(val[0], scale); + ret <<= 4; + ret |= tmp; + return ret; +#endif +} + +/** + * @brief Convert packed fp16x2 to fp4x2. + * + * @param val + * @param scale + * @return __amd_fp4x2_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp4x2_storage_t __amd_cvt_fp16x2_to_fp4x2_scale( + const __amd_fp16x2_storage_t val, const __amd_fp4_interpretation_t, const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(__amd_fp4x2_storage_t[4]) == sizeof(unsigned int)); + union { + unsigned int ui32; + __amd_fp4x2_storage_t fp4x2[4]; + } u{0}; + u.ui32 = __builtin_amdgcn_cvt_scalef32_pk_fp4_f16(u.ui32, val, __amd_scale_to_float(scale), 0); + return u.fp4x2[0]; +#else + using namespace fcbx; + __amd_fp4x2_storage_t ret; + ret = from_float<__amd_fp16_storage_t, Encoding::E2M1, true>(val[1], scale); + ret <<= 4; + ret |= from_float<__amd_fp16_storage_t, Encoding::E2M1, true>(val[0], scale); + return ret; +#endif +} + +/** + * @brief Convert packed fp16x8 to fp4x8. + * + * @param val + * @param scale + * @return __amd_fp4x8_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp4x8_storage_t __amd_cvt_fp16x8_to_fp4x8_scale( + const __amd_fp16x8_storage_t val, const __amd_fp4_interpretation_t, const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(__amd_fp4x2_storage_t[4]) == sizeof(unsigned int)); + static_assert(sizeof(__amd_fp4x2_storage_t[4]) == sizeof(__amd_fp4x8_storage_t)); + union hold_u { + unsigned int ui32; + __amd_fp4x2_storage_t fp4x2[4]; + __amd_fp4x8_storage_t fp4x8; + } ret{0}, tmp{0}; + __amd_fp16x2_storage_t tmp_in{val[0], val[1]}; + tmp.ui32 = + __builtin_amdgcn_cvt_scalef32_pk_fp4_f16(tmp.ui32, tmp_in, __amd_scale_to_float(scale), 0); + ret.fp4x2[0] = tmp.fp4x2[0]; + tmp_in[0] = val[2]; + tmp_in[1] = val[3]; + tmp.ui32 = + __builtin_amdgcn_cvt_scalef32_pk_fp4_f16(tmp.ui32, tmp_in, __amd_scale_to_float(scale), 0); + ret.fp4x2[1] = tmp.fp4x2[0]; + tmp_in[0] = val[4]; + tmp_in[1] = val[5]; + tmp.ui32 = + __builtin_amdgcn_cvt_scalef32_pk_fp4_f16(tmp.ui32, tmp_in, __amd_scale_to_float(scale), 0); + ret.fp4x2[2] = tmp.fp4x2[0]; + tmp_in[0] = val[6]; + tmp_in[1] = val[7]; + tmp.ui32 = + __builtin_amdgcn_cvt_scalef32_pk_fp4_f16(tmp.ui32, tmp_in, __amd_scale_to_float(scale), 0); + ret.fp4x2[3] = tmp.fp4x2[0]; + return ret.fp4x8; +#else + __amd_fp4x8_storage_t ret = 0; + using namespace fcbx; + auto tmp = from_float<__amd_fp16_storage_t, Encoding::E2M1, true>(val[7], scale); + ret <<= 4; + ret |= tmp; + tmp = from_float<__amd_fp16_storage_t, Encoding::E2M1, true>(val[6], scale); + ret <<= 4; + ret |= tmp; + tmp = from_float<__amd_fp16_storage_t, Encoding::E2M1, true>(val[5], scale); + ret <<= 4; + ret |= tmp; + tmp = from_float<__amd_fp16_storage_t, Encoding::E2M1, true>(val[4], scale); + ret <<= 4; + ret |= tmp; + tmp = from_float<__amd_fp16_storage_t, Encoding::E2M1, true>(val[3], scale); + ret <<= 4; + ret |= tmp; + tmp = from_float<__amd_fp16_storage_t, Encoding::E2M1, true>(val[2], scale); + ret <<= 4; + ret |= tmp; + tmp = from_float<__amd_fp16_storage_t, Encoding::E2M1, true>(val[1], scale); + ret <<= 4; + ret |= tmp; + tmp = from_float<__amd_fp16_storage_t, Encoding::E2M1, true>(val[0], scale); + ret <<= 4; + ret |= tmp; + return ret; +#endif +} + +/** + * @brief Convert floatx8 to fp4x8 with stochastic rounding and scale. + * + * @param val + * @param seed + * @param scale + * @return __amd_fp4x8_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp4x8_storage_t __amd_cvt_floatx8_to_fp4x8_sr_scale( + const __amd_floatx8_storage_t val, const __amd_fp4_interpretation_t, const unsigned int seed, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(__amd_fp4x2_storage_t[4]) == sizeof(unsigned int)); + union { + unsigned int ui32; + __amd_fp4x2_storage_t fp4x2[4]; + } tmp{0}; + union { + __amd_fp4x2_storage_t fp4x2[4]; + __amd_fp4x8_storage_t fp4x8; + } ret; + __amd_floatx2_storage_t in{val[0], val[1]}; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_pk_fp4_f32(tmp.ui32, in, seed, + __amd_scale_to_float(scale), 1); + ret.fp4x2[0] = tmp.fp4x2[1]; + in[0] = val[2]; + in[1] = val[3]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_pk_fp4_f32(tmp.ui32, in, seed, + __amd_scale_to_float(scale), 1); + ret.fp4x2[1] = tmp.fp4x2[1]; + in[0] = val[4]; + in[1] = val[5]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_pk_fp4_f32(tmp.ui32, in, seed, + __amd_scale_to_float(scale), 1); + ret.fp4x2[2] = tmp.fp4x2[1]; + in[0] = val[6]; + in[1] = val[7]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_pk_fp4_f32(tmp.ui32, in, seed, + __amd_scale_to_float(scale), 1); + ret.fp4x2[3] = tmp.fp4x2[1]; + return ret.fp4x8; +#else + __amd_fp4x8_storage_t ret = 0; + using namespace fcbx; + auto tmp = from_float(val[7], scale); + ret <<= 4; + ret |= tmp; + tmp = from_float(val[6], scale); + ret <<= 4; + ret |= tmp; + tmp = from_float(val[5], scale); + ret <<= 4; + ret |= tmp; + tmp = from_float(val[4], scale); + ret <<= 4; + ret |= tmp; + tmp = from_float(val[3], scale); + ret <<= 4; + ret |= tmp; + tmp = from_float(val[2], scale); + ret <<= 4; + ret |= tmp; + tmp = from_float(val[1], scale); + ret <<= 4; + ret |= tmp; + tmp = from_float(val[0], scale); + ret <<= 4; + ret |= tmp; + return ret; +#endif +} + +/** + * @brief Convert packed bf16x2 to fp4x2 with stochastic rounding and scale. + * + * @param val + * @param seed + * @param scale + * @return __amd_fp4x2_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp4x2_storage_t __amd_cvt_bf16x2_to_fp4x2_sr_scale( + const __amd_bf16x2_storage_t val, const __amd_fp4_interpretation_t, const unsigned int seed, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(__amd_fp4x2_storage_t[4]) == sizeof(unsigned int)); + union { + unsigned int ui32; + __amd_fp4x2_storage_t fp4x2[4]; + } u{0}; + u.ui32 = __builtin_amdgcn_cvt_scalef32_sr_pk_fp4_bf16(u.ui32, val, seed, + __amd_scale_to_float(scale), 1); + return u.fp4x2[1]; +#else + __amd_fp4x2_storage_t ret; + using namespace fcbx; + ret = from_float_sr<__amd_bf16_storage_t, Encoding::E2M1, true>(val[1], seed, scale); + ret <<= 4; + ret |= from_float_sr<__amd_bf16_storage_t, Encoding::E2M1, true>(val[0], seed, scale); + return ret; +#endif +} + +/** + * @brief Convert bf16x8 to fp4x8 with stochastic rounding and scale. + * + * @param val + * @param seed + * @param scale + * @return __amd_fp4x8_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp4x8_storage_t __amd_cvt_bf16x8_to_fp4x8_sr_scale( + const __amd_bf16x8_storage_t val, const __amd_fp4_interpretation_t, const unsigned int seed, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(__amd_fp4x2_storage_t[4]) == sizeof(unsigned int)); + union { + unsigned int ui32; + __amd_fp4x2_storage_t fp4x2[4]; + } tmp{0}; + union { + __amd_fp4x2_storage_t fp4x2[4]; + __amd_fp4x8_storage_t fp4x8; + } ret; + __amd_bf16x2_storage_t in{val[0], val[1]}; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_pk_fp4_bf16(tmp.ui32, in, seed, + __amd_scale_to_float(scale), 1); + ret.fp4x2[0] = tmp.fp4x2[1]; + in[0] = val[2]; + in[1] = val[3]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_pk_fp4_bf16(tmp.ui32, in, seed, + __amd_scale_to_float(scale), 1); + ret.fp4x2[1] = tmp.fp4x2[1]; + in[0] = val[4]; + in[1] = val[5]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_pk_fp4_bf16(tmp.ui32, in, seed, + __amd_scale_to_float(scale), 1); + ret.fp4x2[2] = tmp.fp4x2[1]; + in[0] = val[6]; + in[1] = val[7]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_pk_fp4_bf16(tmp.ui32, in, seed, + __amd_scale_to_float(scale), 1); + ret.fp4x2[3] = tmp.fp4x2[1]; + return ret.fp4x8; +#else + __amd_fp4x8_storage_t ret = 0; + using namespace fcbx; + auto tmp = from_float_sr<__amd_bf16_storage_t, Encoding::E2M1, true>(val[7], seed, scale); + ret <<= 4; + ret |= tmp; + tmp = from_float_sr<__amd_bf16_storage_t, Encoding::E2M1, true>(val[6], seed, scale); + ret <<= 4; + ret |= tmp; + tmp = from_float_sr<__amd_bf16_storage_t, Encoding::E2M1, true>(val[5], seed, scale); + ret <<= 4; + ret |= tmp; + tmp = from_float_sr<__amd_bf16_storage_t, Encoding::E2M1, true>(val[4], seed, scale); + ret <<= 4; + ret |= tmp; + tmp = from_float_sr<__amd_bf16_storage_t, Encoding::E2M1, true>(val[3], seed, scale); + ret <<= 4; + ret |= tmp; + tmp = from_float_sr<__amd_bf16_storage_t, Encoding::E2M1, true>(val[2], seed, scale); + ret <<= 4; + ret |= tmp; + tmp = from_float_sr<__amd_bf16_storage_t, Encoding::E2M1, true>(val[1], seed, scale); + ret <<= 4; + ret |= tmp; + tmp = from_float_sr<__amd_bf16_storage_t, Encoding::E2M1, true>(val[0], seed, scale); + ret <<= 4; + ret |= tmp; + return ret; +#endif +} + +/** + * @brief Convert packed fp16x2 to fp4x2 with stochastic rounding and scale. + * + * @param val + * @param seed + * @param scale + * @return __amd_fp4x2_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp4x2_storage_t __amd_cvt_fp16x2_to_fp4x2_sr_scale( + const __amd_fp16x2_storage_t val, const __amd_fp4_interpretation_t, const unsigned int seed, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(__amd_fp4x2_storage_t[4]) == sizeof(unsigned int)); + union { + unsigned int ui32; + __amd_fp4x2_storage_t fp4x2[4]; + } u{0}; + u.ui32 = __builtin_amdgcn_cvt_scalef32_sr_pk_fp4_f16(u.ui32, val, seed, + __amd_scale_to_float(scale), 1); + return u.fp4x2[1]; +#else + __amd_fp4x2_storage_t ret; + using namespace fcbx; + ret = from_float_sr<__amd_fp16_storage_t, Encoding::E2M1, true>(val[1], seed, scale); + ret <<= 4; + ret |= from_float_sr<__amd_fp16_storage_t, Encoding::E2M1, true>(val[0], seed, scale); + return ret; +#endif +} + +/** + * @brief Convert fp16x8 to fp4x8 with stochastic rounding and scale. + * + * @param val + * @param seed + * @param scale + * @return __amd_fp4x8_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp4x8_storage_t __amd_cvt_fp16x8_to_fp4x8_sr_scale( + const __amd_fp16x8_storage_t val, const __amd_fp4_interpretation_t, const unsigned int seed, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(__amd_fp4x2_storage_t[4]) == sizeof(unsigned int)); + union { + unsigned int ui32; + __amd_fp4x2_storage_t fp4x2[4]; + } tmp{0}; + union { + __amd_fp4x2_storage_t fp4x2[4]; + __amd_fp4x8_storage_t fp4x8; + } ret; + __amd_fp16x2_storage_t in{val[0], val[1]}; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_pk_fp4_f16(tmp.ui32, in, seed, + __amd_scale_to_float(scale), 1); + ret.fp4x2[0] = tmp.fp4x2[1]; + in[0] = val[2]; + in[1] = val[3]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_pk_fp4_f16(tmp.ui32, in, seed, + __amd_scale_to_float(scale), 1); + ret.fp4x2[1] = tmp.fp4x2[1]; + in[0] = val[4]; + in[1] = val[5]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_pk_fp4_f16(tmp.ui32, in, seed, + __amd_scale_to_float(scale), 1); + ret.fp4x2[2] = tmp.fp4x2[1]; + in[0] = val[6]; + in[1] = val[7]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_pk_fp4_f16(tmp.ui32, in, seed, + __amd_scale_to_float(scale), 1); + ret.fp4x2[3] = tmp.fp4x2[1]; + return ret.fp4x8; +#else + __amd_fp4x8_storage_t ret = 0; + using namespace fcbx; + auto tmp = from_float_sr<__amd_fp16_storage_t, Encoding::E2M1, true>(val[7], seed, scale); + ret <<= 4; + ret |= tmp; + tmp = from_float_sr<__amd_fp16_storage_t, Encoding::E2M1, true>(val[6], seed, scale); + ret <<= 4; + ret |= tmp; + tmp = from_float_sr<__amd_fp16_storage_t, Encoding::E2M1, true>(val[5], seed, scale); + ret <<= 4; + ret |= tmp; + tmp = from_float_sr<__amd_fp16_storage_t, Encoding::E2M1, true>(val[4], seed, scale); + ret <<= 4; + ret |= tmp; + tmp = from_float_sr<__amd_fp16_storage_t, Encoding::E2M1, true>(val[3], seed, scale); + ret <<= 4; + ret |= tmp; + tmp = from_float_sr<__amd_fp16_storage_t, Encoding::E2M1, true>(val[2], seed, scale); + ret <<= 4; + ret |= tmp; + tmp = from_float_sr<__amd_fp16_storage_t, Encoding::E2M1, true>(val[1], seed, scale); + ret <<= 4; + ret |= tmp; + tmp = from_float_sr<__amd_fp16_storage_t, Encoding::E2M1, true>(val[0], seed, scale); + ret <<= 4; + ret |= tmp; + return ret; +#endif +} + +/** + * @brief Convert packed floatx8 to fp8x8 with stochastic rounding and scale. + * + * @param val + * @param interpret + * @param seed + * @param scale + * @return __amd_fp8x8_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp8x8_storage_t __amd_cvt_floatx8_to_fp8x8_sr_scale( + const __amd_floatx8_storage_t val, const __amd_fp8_interpretation_t interpret, + const unsigned int seed, const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(__amd_fp8_storage_t[4]) == sizeof(unsigned int)); + __amd_fp8x8_storage_t ret; + union hold_u { + unsigned int ui32; + __amd_fp8_storage_t fp8[4]; + } tmp{0}; + if (interpret == __AMD_OCP_E4M3) { + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_fp8_f32(tmp.ui32, val[0], seed, + __amd_scale_to_float(scale), 0); + ret[0] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_fp8_f32(tmp.ui32, val[1], seed, + __amd_scale_to_float(scale), 0); + ret[1] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_fp8_f32(tmp.ui32, val[2], seed, + __amd_scale_to_float(scale), 0); + ret[2] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_fp8_f32(tmp.ui32, val[3], seed, + __amd_scale_to_float(scale), 0); + ret[3] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_fp8_f32(tmp.ui32, val[4], seed, + __amd_scale_to_float(scale), 0); + ret[4] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_fp8_f32(tmp.ui32, val[5], seed, + __amd_scale_to_float(scale), 0); + ret[5] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_fp8_f32(tmp.ui32, val[6], seed, + __amd_scale_to_float(scale), 0); + ret[6] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_fp8_f32(tmp.ui32, val[7], seed, + __amd_scale_to_float(scale), 0); + ret[7] = tmp.fp8[0]; + } else { + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_bf8_f32(tmp.ui32, val[0], seed, + __amd_scale_to_float(scale), 0); + ret[0] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_bf8_f32(tmp.ui32, val[1], seed, + __amd_scale_to_float(scale), 0); + ret[1] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_bf8_f32(tmp.ui32, val[2], seed, + __amd_scale_to_float(scale), 0); + ret[2] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_bf8_f32(tmp.ui32, val[3], seed, + __amd_scale_to_float(scale), 0); + ret[3] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_bf8_f32(tmp.ui32, val[4], seed, + __amd_scale_to_float(scale), 0); + ret[4] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_bf8_f32(tmp.ui32, val[5], seed, + __amd_scale_to_float(scale), 0); + ret[5] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_bf8_f32(tmp.ui32, val[6], seed, + __amd_scale_to_float(scale), 0); + ret[6] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_bf8_f32(tmp.ui32, val[7], seed, + __amd_scale_to_float(scale), 0); + ret[7] = tmp.fp8[0]; + } + return ret; +#else + using namespace fcbx; + __amd_fp8x8_storage_t ret; + if (interpret == __AMD_OCP_E4M3) { + ret[0] = from_float_sr(val[0], seed, scale); + ret[1] = from_float_sr(val[1], seed, scale); + ret[2] = from_float_sr(val[2], seed, scale); + ret[3] = from_float_sr(val[3], seed, scale); + ret[4] = from_float_sr(val[4], seed, scale); + ret[5] = from_float_sr(val[5], seed, scale); + ret[6] = from_float_sr(val[6], seed, scale); + ret[7] = from_float_sr(val[7], seed, scale); + } else { + ret[0] = from_float_sr(val[0], seed, scale); + ret[1] = from_float_sr(val[1], seed, scale); + ret[2] = from_float_sr(val[2], seed, scale); + ret[3] = from_float_sr(val[3], seed, scale); + ret[4] = from_float_sr(val[4], seed, scale); + ret[5] = from_float_sr(val[5], seed, scale); + ret[6] = from_float_sr(val[6], seed, scale); + ret[7] = from_float_sr(val[7], seed, scale); + } + return ret; +#endif +} + +/** + * @brief Convret fp16 to fp8 with stochastic rounding and scale. + * + * @param val + * @param interpret + * @param seed + * @param scale + * @return __amd_fp8_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp8_storage_t __amd_cvt_fp16_to_fp8_sr_scale( + const __amd_fp16_storage_t val, const __amd_fp8_interpretation_t interpret, + const unsigned int seed, const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(__amd_fp8_storage_t[4]) == sizeof(unsigned int)); + union u { + unsigned int ui32; + __amd_fp8_storage_t fp8[4]; + } u{0}; + if (interpret == __AMD_OCP_E4M3) { + u.ui32 = + __builtin_amdgcn_cvt_scalef32_sr_fp8_f16(u.ui32, val, seed, __amd_scale_to_float(scale), 0); + } else { + u.ui32 = + __builtin_amdgcn_cvt_scalef32_sr_bf8_f16(u.ui32, val, seed, __amd_scale_to_float(scale), 0); + } + return u.fp8[0]; +#else + using namespace fcbx; + if (interpret == __AMD_OCP_E4M3) { + return from_float_sr<__amd_fp16_storage_t, Encoding::E4M3, true>(val, seed, scale); + } else { + return from_float_sr<__amd_fp16_storage_t, Encoding::E5M2, true>(val, seed, scale); + } +#endif +} + +/** + * @brief Convert packed fp16x8 to fp8x8 with stochastic rounding and scale. + * + * @param val + * @param interpret + * @param seed + * @param scale + * @return __amd_fp8x8_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp8x8_storage_t __amd_cvt_fp16x8_to_fp8x8_sr_scale( + const __amd_fp16x8_storage_t val, const __amd_fp8_interpretation_t interpret, + const unsigned int seed, const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(__amd_fp8_storage_t[4]) == sizeof(unsigned int)); + __amd_fp8x8_storage_t ret; + union hold_u { + unsigned int ui32; + __amd_fp8_storage_t fp8[4]; + } tmp{0}; + if (interpret == __AMD_OCP_E4M3) { + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_fp8_f16(tmp.ui32, val[0], seed, + __amd_scale_to_float(scale), 0); + ret[0] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_fp8_f16(tmp.ui32, val[1], seed, + __amd_scale_to_float(scale), 0); + ret[1] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_fp8_f16(tmp.ui32, val[2], seed, + __amd_scale_to_float(scale), 0); + ret[2] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_fp8_f16(tmp.ui32, val[3], seed, + __amd_scale_to_float(scale), 0); + ret[3] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_fp8_f16(tmp.ui32, val[4], seed, + __amd_scale_to_float(scale), 0); + ret[4] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_fp8_f16(tmp.ui32, val[5], seed, + __amd_scale_to_float(scale), 0); + ret[5] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_fp8_f16(tmp.ui32, val[6], seed, + __amd_scale_to_float(scale), 0); + ret[6] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_fp8_f16(tmp.ui32, val[7], seed, + __amd_scale_to_float(scale), 0); + ret[7] = tmp.fp8[0]; + } else { + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_bf8_f16(tmp.ui32, val[0], seed, + __amd_scale_to_float(scale), 0); + ret[0] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_bf8_f16(tmp.ui32, val[1], seed, + __amd_scale_to_float(scale), 0); + ret[1] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_bf8_f16(tmp.ui32, val[2], seed, + __amd_scale_to_float(scale), 0); + ret[2] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_bf8_f16(tmp.ui32, val[3], seed, + __amd_scale_to_float(scale), 0); + ret[3] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_bf8_f16(tmp.ui32, val[4], seed, + __amd_scale_to_float(scale), 0); + ret[4] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_bf8_f16(tmp.ui32, val[5], seed, + __amd_scale_to_float(scale), 0); + ret[5] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_bf8_f16(tmp.ui32, val[6], seed, + __amd_scale_to_float(scale), 0); + ret[6] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_bf8_f16(tmp.ui32, val[7], seed, + __amd_scale_to_float(scale), 0); + ret[7] = tmp.fp8[0]; + } + return ret; +#else + using namespace fcbx; + __amd_fp8x8_storage_t ret; + if (interpret == __AMD_OCP_E4M3) { + ret[0] = from_float_sr<__amd_fp16_storage_t, Encoding::E4M3, true>(val[0], seed, scale); + ret[1] = from_float_sr<__amd_fp16_storage_t, Encoding::E4M3, true>(val[1], seed, scale); + ret[2] = from_float_sr<__amd_fp16_storage_t, Encoding::E4M3, true>(val[2], seed, scale); + ret[3] = from_float_sr<__amd_fp16_storage_t, Encoding::E4M3, true>(val[3], seed, scale); + ret[4] = from_float_sr<__amd_fp16_storage_t, Encoding::E4M3, true>(val[4], seed, scale); + ret[5] = from_float_sr<__amd_fp16_storage_t, Encoding::E4M3, true>(val[5], seed, scale); + ret[6] = from_float_sr<__amd_fp16_storage_t, Encoding::E4M3, true>(val[6], seed, scale); + ret[7] = from_float_sr<__amd_fp16_storage_t, Encoding::E4M3, true>(val[7], seed, scale); + } else { + ret[0] = from_float_sr<__amd_fp16_storage_t, Encoding::E5M2, true>(val[0], seed, scale); + ret[1] = from_float_sr<__amd_fp16_storage_t, Encoding::E5M2, true>(val[1], seed, scale); + ret[2] = from_float_sr<__amd_fp16_storage_t, Encoding::E5M2, true>(val[2], seed, scale); + ret[3] = from_float_sr<__amd_fp16_storage_t, Encoding::E5M2, true>(val[3], seed, scale); + ret[4] = from_float_sr<__amd_fp16_storage_t, Encoding::E5M2, true>(val[4], seed, scale); + ret[5] = from_float_sr<__amd_fp16_storage_t, Encoding::E5M2, true>(val[5], seed, scale); + ret[6] = from_float_sr<__amd_fp16_storage_t, Encoding::E5M2, true>(val[6], seed, scale); + ret[7] = from_float_sr<__amd_fp16_storage_t, Encoding::E5M2, true>(val[7], seed, scale); + } + return ret; +#endif +} + +/** + * @brief Convert bf16 to fp8 with stochastic rounding and scale + * + * @param val + * @param interpret + * @param seed + * @param scale + * @return __amd_fp8_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp8_storage_t __amd_cvt_bf16_to_fp8_sr_scale( + const __amd_bf16_storage_t val, const __amd_fp8_interpretation_t interpret, + const unsigned int seed, const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(__amd_fp8_storage_t[4]) == sizeof(unsigned int)); + union u { + unsigned int ui32; + __amd_fp8_storage_t fp8[4]; + } u{0}; + if (interpret == __AMD_OCP_E4M3) { + u.ui32 = __builtin_amdgcn_cvt_scalef32_sr_fp8_bf16(u.ui32, val, seed, + __amd_scale_to_float(scale), 0); + } else { + u.ui32 = __builtin_amdgcn_cvt_scalef32_sr_bf8_bf16(u.ui32, val, seed, + __amd_scale_to_float(scale), 0); + } + return u.fp8[0]; +#else + using namespace fcbx; + if (interpret == __AMD_OCP_E4M3) { + return from_float_sr<__amd_bf16_storage_t, Encoding::E4M3, true>(val, seed, scale); + } else { + return from_float_sr<__amd_bf16_storage_t, Encoding::E5M2, true>(val, seed, scale); + } +#endif +} + +/** + * @brief Convert packed bf16x8 to fp8x8 with stochastic rounding and scale. + * + * @param val + * @param interpret + * @param seed + * @param scale + * @return __amd_fp8x8_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp8x8_storage_t __amd_cvt_bf16x8_to_fp8x8_sr_scale( + const __amd_bf16x8_storage_t val, const __amd_fp8_interpretation_t interpret, + const unsigned int seed, const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(__amd_fp8_storage_t[4]) == sizeof(unsigned int)); + __amd_fp8x8_storage_t ret; + union hold_u { + unsigned int ui32; + __amd_fp8_storage_t fp8[4]; + } tmp{0}; + if (interpret == __AMD_OCP_E4M3) { + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_fp8_bf16(/*tmp.ui32*/ 0, val[0], seed, + __amd_scale_to_float(scale), 0); + ret[0] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_fp8_bf16(/*tmp.ui32*/ 0, val[1], seed, + __amd_scale_to_float(scale), 0); + ret[1] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_fp8_bf16(/*tmp.ui32*/ 0, val[2], seed, + __amd_scale_to_float(scale), 0); + ret[2] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_fp8_bf16(/*tmp.ui32*/ 0, val[3], seed, + __amd_scale_to_float(scale), 0); + ret[3] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_fp8_bf16(/*tmp.ui32*/ 0, val[4], seed, + __amd_scale_to_float(scale), 0); + ret[4] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_fp8_bf16(/*tmp.ui32*/ 0, val[5], seed, + __amd_scale_to_float(scale), 0); + ret[5] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_fp8_bf16(/*tmp.ui32*/ 0, val[6], seed, + __amd_scale_to_float(scale), 0); + ret[6] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_fp8_bf16(/*tmp.ui32*/ 0, val[7], seed, + __amd_scale_to_float(scale), 0); + ret[7] = tmp.fp8[0]; + } else { + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_bf8_bf16(/*tmp.ui32*/ 0, val[0], seed, + __amd_scale_to_float(scale), 0); + ret[0] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_bf8_bf16(/*tmp.ui32*/ 0, val[1], seed, + __amd_scale_to_float(scale), 0); + ret[1] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_bf8_bf16(/*tmp.ui32*/ 0, val[2], seed, + __amd_scale_to_float(scale), 0); + ret[2] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_bf8_bf16(/*tmp.ui32*/ 0, val[3], seed, + __amd_scale_to_float(scale), 0); + ret[3] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_bf8_bf16(/*tmp.ui32*/ 0, val[4], seed, + __amd_scale_to_float(scale), 0); + ret[4] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_bf8_bf16(/*tmp.ui32*/ 0, val[5], seed, + __amd_scale_to_float(scale), 0); + ret[5] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_bf8_bf16(/*tmp.ui32*/ 0, val[6], seed, + __amd_scale_to_float(scale), 0); + ret[6] = tmp.fp8[0]; + tmp.ui32 = __builtin_amdgcn_cvt_scalef32_sr_bf8_bf16(/*tmp.ui32*/ 0, val[7], seed, + __amd_scale_to_float(scale), 0); + ret[7] = tmp.fp8[0]; + } + return ret; +#else + using namespace fcbx; + __amd_fp8x8_storage_t ret; + if (interpret == __AMD_OCP_E4M3) { + ret[0] = from_float_sr<__amd_bf16_storage_t, Encoding::E4M3, true>(val[0], seed, scale); + ret[1] = from_float_sr<__amd_bf16_storage_t, Encoding::E4M3, true>(val[1], seed, scale); + ret[2] = from_float_sr<__amd_bf16_storage_t, Encoding::E4M3, true>(val[2], seed, scale); + ret[3] = from_float_sr<__amd_bf16_storage_t, Encoding::E4M3, true>(val[3], seed, scale); + ret[4] = from_float_sr<__amd_bf16_storage_t, Encoding::E4M3, true>(val[4], seed, scale); + ret[5] = from_float_sr<__amd_bf16_storage_t, Encoding::E4M3, true>(val[5], seed, scale); + ret[6] = from_float_sr<__amd_bf16_storage_t, Encoding::E4M3, true>(val[6], seed, scale); + ret[7] = from_float_sr<__amd_bf16_storage_t, Encoding::E4M3, true>(val[7], seed, scale); + } else { + ret[0] = from_float_sr<__amd_bf16_storage_t, Encoding::E5M2, true>(val[0], seed, scale); + ret[1] = from_float_sr<__amd_bf16_storage_t, Encoding::E5M2, true>(val[1], seed, scale); + ret[2] = from_float_sr<__amd_bf16_storage_t, Encoding::E5M2, true>(val[2], seed, scale); + ret[3] = from_float_sr<__amd_bf16_storage_t, Encoding::E5M2, true>(val[3], seed, scale); + ret[4] = from_float_sr<__amd_bf16_storage_t, Encoding::E5M2, true>(val[4], seed, scale); + ret[5] = from_float_sr<__amd_bf16_storage_t, Encoding::E5M2, true>(val[5], seed, scale); + ret[6] = from_float_sr<__amd_bf16_storage_t, Encoding::E5M2, true>(val[6], seed, scale); + ret[7] = from_float_sr<__amd_bf16_storage_t, Encoding::E5M2, true>(val[7], seed, scale); + } + return ret; +#endif +} + +/** + * @brief Convert fp8 to fp16. + * + * @param val fp8 number + * @param interpret interpretation of fp8 + * @return __amd_fp16_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp16_storage_t +__amd_cvt_fp8_to_fp16(const __amd_fp8_storage_t val, const __amd_fp8_interpretation_t interpret) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + __amd_fp16x2_storage_t ret; + if (interpret == __AMD_OCP_E4M3) { + ret = __builtin_amdgcn_cvt_scalef32_f16_fp8(ret, val, __amd_scale_to_float(0), 0, false); + } else { + ret = __builtin_amdgcn_cvt_scalef32_f16_bf8(ret, val, __amd_scale_to_float(0), 0, false); + } + return ret[0]; +#else + using namespace fcbx; + if (interpret == __AMD_OCP_E4M3) { + return to_float<__amd_fp16_storage_t, Encoding::E4M3, true>(val, 0); + } else { + return to_float<__amd_fp16_storage_t, Encoding::E5M2, true>(val, 0); + } +#endif +} + +/** + * @brief Convert packed fp16 vector of size 2 to fp8 vector of size 2 + * + * @param val fp8x2 value + * @param interpret interpretation of fp8 + * @return __amd_fp16x2_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp16x2_storage_t __amd_cvt_fp8x2_to_fp16x2( + const __amd_fp8x2_storage_t val, const __amd_fp8_interpretation_t interpret) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(unsigned int) == sizeof(__amd_fp8x2_storage_t[2])); + union { + __amd_fp8x2_storage_t fp8x2[2]; + unsigned int ui32; + } u; + u.fp8x2[0] = val; + return interpret == __AMD_OCP_E4M3 + ? __builtin_amdgcn_cvt_scalef32_pk_f16_fp8(u.ui32, __amd_scale_to_float(0), false) + : __builtin_amdgcn_cvt_scalef32_pk_f16_bf8(u.ui32, __amd_scale_to_float(0), false); +#else + using namespace fcbx; + __amd_fp16x2_storage_t ret; + if (interpret == __AMD_OCP_E4M3) { + ret[0] = to_float<__amd_fp16_storage_t, Encoding::E4M3, true>(val & 0xFF, 0); + ret[1] = to_float<__amd_fp16_storage_t, Encoding::E4M3, true>(val >> 8, 0); + } else { + ret[0] = to_float<__amd_fp16_storage_t, Encoding::E5M2, true>(val & 0xFF, 0); + ret[1] = to_float<__amd_fp16_storage_t, Encoding::E5M2, true>(val >> 8, 0); + } + return ret; +#endif +} + +/** + * @brief Convert packed fp16 of size 2 to fp8 + * + * @param val fp8x2 value + * @param interpret interpretation of fp8 + * @return __amd_fp8x2_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp8x2_storage_t __amd_cvt_fp16x2_to_fp8x2( + const __amd_fp16x2_storage_t val, const __amd_fp8_interpretation_t interpret) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(__amd_shortx2_storage_t) == sizeof(__amd_fp8x2_storage_t[2])); + union { + __amd_shortx2_storage_t shortx2; + __amd_fp8x2_storage_t fp8x2[2]; + } u{0}; + u.shortx2 = interpret == __AMD_OCP_E4M3 + ? __builtin_amdgcn_cvt_scalef32_pk_fp8_f16(u.shortx2, val, __amd_scale_to_float(0), false) + : __builtin_amdgcn_cvt_scalef32_pk_bf8_f16(u.shortx2, val, __amd_scale_to_float(0), false); + return u.fp8x2[0]; +#else + using namespace fcbx; + __amd_fp8x2_storage_t ret; + if (interpret == __AMD_OCP_E4M3) { + ret = from_float<__amd_fp16_storage_t, Encoding::E4M3, true>(val[1], 0); + ret <<= 8; + ret |= from_float<__amd_fp16_storage_t, Encoding::E4M3, true>(val[0], 0); + } else { + ret = from_float<__amd_fp16_storage_t, Encoding::E5M2, true>(val[1], 0); + ret <<= 8; + ret |= from_float<__amd_fp16_storage_t, Encoding::E5M2, true>(val[0], 0); + } + return ret; +#endif +} + +/** + * @brief Convert fp16 pack 8 to fp8 packed 8 + * + * @param val fp16x8 value + * @param interpret interpretation of fp8 + * @param scale + * @return __amd_fp8x8_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp8x8_storage_t __amd_cvt_fp16x8_to_fp8x8_scale( + const __amd_fp16x8_storage_t val, const __amd_fp8_interpretation_t interpret, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(__amd_shortx2_storage_t) == sizeof(__amd_fp8x2_storage_t[2])); + static_assert(sizeof(__amd_fp8x8_storage_t) == sizeof(__amd_fp8x2_storage_t[4])); + union { + __amd_shortx2_storage_t shortx2; + __amd_fp8x2_storage_t fp8x2[2]; + } tmp{0}; + union { + __amd_fp8x2_storage_t fp8x2[4]; + __amd_fp8x8_storage_t fp8x8; + } ret; + __amd_fp16x2_storage_t in{val[0], val[1]}; + if (interpret == __AMD_OCP_E4M3) { + tmp.shortx2 = __builtin_amdgcn_cvt_scalef32_pk_fp8_f16(tmp.shortx2, in, + __amd_scale_to_float(scale), false); + ret.fp8x2[0] = tmp.fp8x2[0]; + in[0] = val[2]; + in[1] = val[3]; + tmp.shortx2 = __builtin_amdgcn_cvt_scalef32_pk_fp8_f16(tmp.shortx2, in, + __amd_scale_to_float(scale), false); + ret.fp8x2[1] = tmp.fp8x2[0]; + in[0] = val[4]; + in[1] = val[5]; + tmp.shortx2 = __builtin_amdgcn_cvt_scalef32_pk_fp8_f16(tmp.shortx2, in, + __amd_scale_to_float(scale), false); + ret.fp8x2[2] = tmp.fp8x2[0]; + in[0] = val[6]; + in[1] = val[7]; + tmp.shortx2 = __builtin_amdgcn_cvt_scalef32_pk_fp8_f16(tmp.shortx2, in, + __amd_scale_to_float(scale), false); + ret.fp8x2[3] = tmp.fp8x2[0]; + } else { + tmp.shortx2 = __builtin_amdgcn_cvt_scalef32_pk_bf8_f16(tmp.shortx2, in, + __amd_scale_to_float(scale), false); + ret.fp8x2[0] = tmp.fp8x2[0]; + in[0] = val[2]; + in[1] = val[3]; + tmp.shortx2 = __builtin_amdgcn_cvt_scalef32_pk_bf8_f16(tmp.shortx2, in, + __amd_scale_to_float(scale), false); + ret.fp8x2[1] = tmp.fp8x2[0]; + in[0] = val[4]; + in[1] = val[5]; + tmp.shortx2 = __builtin_amdgcn_cvt_scalef32_pk_bf8_f16(tmp.shortx2, in, + __amd_scale_to_float(scale), false); + ret.fp8x2[2] = tmp.fp8x2[0]; + in[0] = val[6]; + in[1] = val[7]; + tmp.shortx2 = __builtin_amdgcn_cvt_scalef32_pk_bf8_f16(tmp.shortx2, in, + __amd_scale_to_float(scale), false); + ret.fp8x2[3] = tmp.fp8x2[0]; + } + return ret.fp8x8; +#else + __amd_fp8x8_storage_t ret; + using namespace fcbx; + if (interpret == __AMD_OCP_E4M3) { + ret[0] = from_float<__amd_fp16_storage_t, Encoding::E4M3, true>(val[0], scale); + ret[1] = from_float<__amd_fp16_storage_t, Encoding::E4M3, true>(val[1], scale); + ret[2] = from_float<__amd_fp16_storage_t, Encoding::E4M3, true>(val[2], scale); + ret[3] = from_float<__amd_fp16_storage_t, Encoding::E4M3, true>(val[3], scale); + ret[4] = from_float<__amd_fp16_storage_t, Encoding::E4M3, true>(val[4], scale); + ret[5] = from_float<__amd_fp16_storage_t, Encoding::E4M3, true>(val[5], scale); + ret[6] = from_float<__amd_fp16_storage_t, Encoding::E4M3, true>(val[6], scale); + ret[7] = from_float<__amd_fp16_storage_t, Encoding::E4M3, true>(val[7], scale); + } else { + ret[0] = from_float<__amd_fp16_storage_t, Encoding::E5M2, true>(val[0], scale); + ret[1] = from_float<__amd_fp16_storage_t, Encoding::E5M2, true>(val[1], scale); + ret[2] = from_float<__amd_fp16_storage_t, Encoding::E5M2, true>(val[2], scale); + ret[3] = from_float<__amd_fp16_storage_t, Encoding::E5M2, true>(val[3], scale); + ret[4] = from_float<__amd_fp16_storage_t, Encoding::E5M2, true>(val[4], scale); + ret[5] = from_float<__amd_fp16_storage_t, Encoding::E5M2, true>(val[5], scale); + ret[6] = from_float<__amd_fp16_storage_t, Encoding::E5M2, true>(val[6], scale); + ret[7] = from_float<__amd_fp16_storage_t, Encoding::E5M2, true>(val[7], scale); + } + return ret; +#endif +} + +/** + * @brief Convert float pack 8 to fp8 packed 8 + * + * @param val floatx8 value + * @param interpret interpretation of fp8 + * @param scale + * @return __amd_fp8x8_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp8x8_storage_t __amd_cvt_floatx8_to_fp8x8_scale( + const __amd_floatx8_storage_t val, const __amd_fp8_interpretation_t interpret, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(__amd_shortx2_storage_t) == sizeof(__amd_fp8x2_storage_t[2])); + static_assert(sizeof(__amd_fp8x8_storage_t) == sizeof(__amd_fp8x2_storage_t[4])); + union { + __amd_shortx2_storage_t shortx2; + __amd_fp8x2_storage_t fp8x2[2]; + } u{0}; + union { + __amd_fp8x2_storage_t fp8x2[4]; + __amd_fp8x8_storage_t fp8x8; + } ret; + if (interpret == __AMD_OCP_E4M3) { + u.shortx2 = __builtin_amdgcn_cvt_scalef32_pk_fp8_f32(u.shortx2, val[0], val[1], + __amd_scale_to_float(scale), false); + ret.fp8x2[0] = u.fp8x2[0]; + u.shortx2 = __builtin_amdgcn_cvt_scalef32_pk_fp8_f32(u.shortx2, val[2], val[3], + __amd_scale_to_float(scale), false); + ret.fp8x2[1] = u.fp8x2[0]; + u.shortx2 = __builtin_amdgcn_cvt_scalef32_pk_fp8_f32(u.shortx2, val[4], val[5], + __amd_scale_to_float(scale), false); + ret.fp8x2[2] = u.fp8x2[0]; + u.shortx2 = __builtin_amdgcn_cvt_scalef32_pk_fp8_f32(u.shortx2, val[6], val[7], + __amd_scale_to_float(scale), false); + ret.fp8x2[3] = u.fp8x2[0]; + } else { + u.shortx2 = __builtin_amdgcn_cvt_scalef32_pk_bf8_f32(u.shortx2, val[0], val[1], + __amd_scale_to_float(scale), false); + ret.fp8x2[0] = u.fp8x2[0]; + u.shortx2 = __builtin_amdgcn_cvt_scalef32_pk_bf8_f32(u.shortx2, val[2], val[3], + __amd_scale_to_float(scale), false); + ret.fp8x2[1] = u.fp8x2[0]; + u.shortx2 = __builtin_amdgcn_cvt_scalef32_pk_bf8_f32(u.shortx2, val[4], val[5], + __amd_scale_to_float(scale), false); + ret.fp8x2[2] = u.fp8x2[0]; + u.shortx2 = __builtin_amdgcn_cvt_scalef32_pk_bf8_f32(u.shortx2, val[6], val[7], + __amd_scale_to_float(scale), false); + ret.fp8x2[3] = u.fp8x2[0]; + } + return ret.fp8x8; +#else + using namespace fcbx; + __amd_fp8x8_storage_t ret; + if (interpret == __AMD_OCP_E4M3) { + ret[0] = from_float(val[0], scale); + ret[1] = from_float(val[1], scale); + ret[2] = from_float(val[2], scale); + ret[3] = from_float(val[3], scale); + ret[4] = from_float(val[4], scale); + ret[5] = from_float(val[5], scale); + ret[6] = from_float(val[6], scale); + ret[7] = from_float(val[7], scale); + } else { + ret[0] = from_float(val[0], scale); + ret[1] = from_float(val[1], scale); + ret[2] = from_float(val[2], scale); + ret[3] = from_float(val[3], scale); + ret[4] = from_float(val[4], scale); + ret[5] = from_float(val[5], scale); + ret[6] = from_float(val[6], scale); + ret[7] = from_float(val[7], scale); + } + return ret; +#endif +} + +/** + * @brief Convert fp16 to fp8 with stochastic rounding. + * + * @param val + * @param interpret + * @param sr + * @return __amd_fp8_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp8_storage_t __amd_cvt_fp16_to_fp8_sr( + const __amd_fp16_storage_t val, const __amd_fp8_interpretation_t interpret, const short sr) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(__amd_fp8_storage_t[4]) == sizeof(unsigned int)); + union u { + unsigned int ui32; + __amd_fp8_storage_t fp8[4]; + } u{0}; + if (interpret == __AMD_OCP_E4M3) { + u.ui32 = __builtin_amdgcn_cvt_scalef32_sr_fp8_f16(u.ui32, val, sr, __amd_scale_to_float(0), 0); + } else { + u.ui32 = __builtin_amdgcn_cvt_scalef32_sr_bf8_f16(u.ui32, val, sr, __amd_scale_to_float(0), 0); + } + return u.fp8[0]; +#else + using namespace fcbx; + return interpret == __AMD_OCP_E4M3 + ? from_float_sr<__amd_fp16_storage_t, Encoding::E4M3, true>(val, sr, 0) + : from_float_sr<__amd_fp16_storage_t, Encoding::E5M2, true>(val, sr, 0); +#endif +} + +/** + * @brief Convert __amd_floatx2_storage_t to hip float2 type + * + * @param val + * @return float2 + */ +__OCP_FP_HOST_DEVICE_STATIC__ float2 +__amd_cvt_floatx2_to_float2(const __amd_floatx2_storage_t val) { + return float2{val[0], val[1]}; +} + +/** + * @brief Convert __amd_fp16_storage_t type to __half type + * + * @param val + * @return __half + */ +__OCP_FP_HOST_DEVICE_STATIC__ __half __amd_cvt_fp16_to_half(const __amd_fp16_storage_t val) { + __half_raw hr; + hr.data = val; + return hr; +} + +/** + * @brief Convert __amd_fp16x2_storage_t to __half2 type + * + * @param val + * @return __half2 + */ +__OCP_FP_HOST_DEVICE_STATIC__ __half2 __amd_cvt_fp16x2_to_half2(const __amd_fp16x2_storage_t val) { + __half2_raw hr; + hr.data = val; + return hr; +} + +/** + * @brief convert __half to __amd_fp16_storage_t + * + * @param val + * @return __amd_fp16_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp16_storage_t __amd_cvt_half_to_fp16(const __half val) { + __half_raw tmp = val; + return tmp.data; +} + +/** + * @brief Convert __half2 to __amd_fp16x2_storage_t + * + * @param val + * @return __amd_fp16x2_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_fp16x2_storage_t __amd_cvt_half2_to_fp16x2(const __half2 val) { + __half2_raw tmp = val; + return tmp.data; +} + +/** + * @brief Convert __amd_bf16_storage_t to __hip_bfloat16 + * + * @param val + * @return __hip_bfloat16 + */ +__OCP_FP_HOST_DEVICE_STATIC__ __hip_bfloat16 +__amd_cvt_bf16_to_hipbf16(const __amd_bf16_storage_t val) { + static_assert(sizeof(__hip_bfloat16) == sizeof(__amd_bf16_storage_t)); + union { + __amd_bf16_storage_t bf16; + __hip_bfloat16 hip_bf16; + } u{val}; + return u.hip_bf16; +} + +/** + * @brief Convert __amd_bf16x2_storage_t to __hip_bfloat162 + * + * @param val + * @return __hip_bfloat162 + */ +__OCP_FP_HOST_DEVICE_STATIC__ __hip_bfloat162 +__amd_cvt_bf16x2_to_hipbf162(const __amd_bf16x2_storage_t val) { + static_assert(sizeof(__hip_bfloat162) == sizeof(__amd_bf16x2_storage_t)); + union { + __amd_bf16x2_storage_t bf16; + __hip_bfloat162 hip_bf16; + } u{val}; + return u.hip_bf16; +} + +/** + * @brief Convert __hip_bfloat16 to __amd_bf16_storage_t + * + * @param val + * @return __amd_bf16_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_bf16_storage_t +__amd_cvt_hipbf16_to_bf16(const __hip_bfloat16 val) { + static_assert(sizeof(__hip_bfloat16) == sizeof(__amd_bf16_storage_t)); + union { + __hip_bfloat16 hip_bf16; + __amd_bf16_storage_t bf16; + } u{val}; + return u.bf16; +} + +/** + * @brief Convert __hip_bfloat162 to __amd_bf16x2_storage_t + * + * @param val + * @return __amd_bf16x2_storage_t + */ +__OCP_FP_HOST_DEVICE_STATIC__ __amd_bf16x2_storage_t +__amd_cvt_hipbf162_to_bf16x2(const __hip_bfloat162 val) { + static_assert(sizeof(__hip_bfloat162) == sizeof(__amd_bf16x2_storage_t)); + union { + __hip_bfloat162 hip_bf16; + __amd_bf16x2_storage_t bf16; + } u{val}; + return u.bf16; +} diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_ocp_fp_cxx.hpp b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_ocp_fp_cxx.hpp new file mode 100644 index 0000000000..b56bd73ac3 --- /dev/null +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_ocp_fp_cxx.hpp @@ -0,0 +1,1056 @@ +#pragma once + +#include "amd_hip_ocp_host.hpp" + +static_assert(sizeof(unsigned int) == sizeof(__amd_fp8_storage_t[4])); +static_assert(sizeof(uint32_t) == sizeof(__amd_fp8_storage_t[4])); +static_assert(sizeof(int) == sizeof(__amd_fp8x2_storage_t[2])); +static_assert(sizeof(uint32_t) == sizeof(__amd_fp8x2_storage_t[2])); +static_assert(sizeof(__amd_shortx2_storage_t) == sizeof(__amd_fp8x2_storage_t[2])); + +struct __hipext_ocp_fp8_e4m3 { + __amd_fp8_storage_t __x; + static const __amd_fp8_interpretation_t __default_interpret = __AMD_OCP_E4M3; + + __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp8_e4m3() = default; + + explicit __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp8_e4m3(const float in) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + union { + unsigned int ui32; + __amd_fp8_storage_t fp8[4]; + } u{0}; + u.ui32 = __builtin_amdgcn_cvt_pk_fp8_f32(in, in, 0, false); + __x = u.fp8[0]; +#else + using namespace fcbx; + union { + uint32_t ui32; + __amd_fp8_storage_t fp8[4]; + } u{0}; + u.ui32 = from_float(in, 0 /*scale*/); + __x = u.fp8[0]; +#endif + } + + explicit __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp8_e4m3(const float in, const unsigned int seed) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + union { + unsigned int ui32; + __amd_fp8_storage_t fp8[4]; + } u{0}; + u.ui32 = __builtin_amdgcn_cvt_sr_fp8_f32(in, seed, 0, 0); + __x = u.fp8[0]; +#else + using namespace fcbx; + union { + uint32_t ui32; + __amd_fp8_storage_t fp8[4]; + } u{0}; + u.ui32 = from_float_sr(in, seed, 0 /*scale*/); + __x = u.fp8[0]; +#endif + } + + explicit __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp8_e4m3(const float in, const unsigned int seed, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + union { + unsigned int ui32; + __amd_fp8_storage_t fp8[4]; + } u{0}; + u.ui32 = + __builtin_amdgcn_cvt_scalef32_sr_fp8_f32(u.ui32, in, seed, __amd_scale_to_float(scale), 0); + __x = u.fp8[0]; +#else + using namespace fcbx; + union { + uint32_t ui32; + __amd_fp8_storage_t fp8[4]; + } u{0}; + u.ui32 = from_float_sr(in, seed, scale); + __x = u.fp8[0]; +#endif + } + + explicit __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp8_e4m3(const __amd_fp16_storage_t in, + const unsigned int seed, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + union u { + unsigned int ui32; + __amd_fp8_storage_t fp8[4]; + } u{0}; + u.ui32 = + __builtin_amdgcn_cvt_scalef32_sr_fp8_f16(u.ui32, in, seed, __amd_scale_to_float(scale), 0); + __x = u.fp8[0]; +#else + using namespace fcbx; + union { + uint32_t ui32; + __amd_fp8_storage_t fp8[4]; + } u{0}; + u.ui32 = from_float_sr<__amd_fp16_storage_t, Encoding::E4M3, true>(in, seed, scale); + __x = u.fp8[0]; +#endif + } + + explicit __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp8_e4m3(const __amd_bf16_storage_t in, + const unsigned int seed, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(__amd_fp8_storage_t[4]) == sizeof(unsigned int)); + union u { + unsigned int ui32; + __amd_fp8_storage_t fp8[4]; + } u{0}; + u.ui32 = + __builtin_amdgcn_cvt_scalef32_sr_fp8_bf16(u.ui32, in, seed, __amd_scale_to_float(scale), 0); + __x = u.fp8[0]; +#else + using namespace fcbx; + union { + uint32_t ui32; + __amd_fp8_storage_t fp8[4]; + } u{0}; + u.ui32 = from_float_sr<__amd_bf16_storage_t, Encoding::E4M3, true>(in, seed, scale); + __x = u.fp8[0]; +#endif + } + + __OCP_FP_HOST_DEVICE__ __amd_fp16_storage_t get_scaled_fp16(const __amd_scale_t scale) const { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + __amd_fp16x2_storage_t ret; + ret = __builtin_amdgcn_cvt_scalef32_f16_fp8(ret, __x, __amd_scale_to_float(scale), 0, false); + return ret[0]; +#else + using namespace fcbx; + return to_float<__amd_fp16_storage_t, Encoding::E4M3, true>(__x, scale); +#endif + } + + __OCP_FP_HOST_DEVICE__ __amd_bf16_storage_t get_scaled_bf16(const __amd_scale_t scale) const { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(unsigned int) == sizeof(__amd_fp8_storage_t[4])); + union { + __amd_fp8_storage_t fp8[4]; + unsigned int ui32; + } u; + u.fp8[0] = __x; + auto ret = + __builtin_amdgcn_cvt_scalef32_pk_bf16_fp8(u.ui32, __amd_scale_to_float(scale), false); + return ret[0]; +#else + using namespace fcbx; + return to_float<__amd_bf16_storage_t, Encoding::E4M3, true>(__x, scale); +#endif + } + + __OCP_FP_HOST_DEVICE__ float get_scaled_float(const __amd_scale_t scale) const { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + return __builtin_amdgcn_cvt_scalef32_f32_fp8(__x, __amd_scale_to_float(scale), 0); +#else + using namespace fcbx; + return to_float(__x, scale); +#endif + } + + __OCP_FP_HOST_DEVICE__ operator float() const { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + return __builtin_amdgcn_cvt_f32_fp8(__x, 0); +#else + using namespace fcbx; + return to_float(__x, 0); +#endif + } +}; + +struct __hipext_ocp_fp8_e5m2 { + __amd_fp8_storage_t __x; + static const __amd_fp8_interpretation_t __default_interpret = __AMD_OCP_E5M2; + + __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp8_e5m2() = default; + + explicit __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp8_e5m2(const float in) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + union { + unsigned int ui32; + __amd_fp8_storage_t fp8[4]; + } u{0}; + u.ui32 = __builtin_amdgcn_cvt_pk_bf8_f32(in, in, 0, false); + __x = u.fp8[0]; +#else + using namespace fcbx; + union { + uint32_t ui32; + __amd_fp8_storage_t fp8[4]; + } u{0}; + u.ui32 = from_float(in, 0 /*scale*/); + __x = u.fp8[0]; +#endif + } + + explicit __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp8_e5m2(const float in, const unsigned int seed) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + union { + unsigned int ui32; + __amd_fp8_storage_t fp8[4]; + } u{0}; + u.ui32 = __builtin_amdgcn_cvt_sr_bf8_f32(in, seed, 0, 0); + __x = u.fp8[0]; +#else + using namespace fcbx; + union { + uint32_t ui32; + __amd_fp8_storage_t fp8[4]; + } u{0}; + u.ui32 = from_float_sr(in, seed, 0 /*scale*/); + __x = u.fp8[0]; +#endif + } + + explicit __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp8_e5m2(const float in, const unsigned int seed, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + union { + unsigned int ui32; + __amd_fp8_storage_t fp8[4]; + } u{0}; + u.ui32 = + __builtin_amdgcn_cvt_scalef32_sr_bf8_f32(u.ui32, in, seed, __amd_scale_to_float(scale), 0); + __x = u.fp8[0]; +#else + using namespace fcbx; + union { + uint32_t ui32; + __amd_fp8_storage_t fp8[4]; + } u{0}; + u.ui32 = from_float_sr(in, seed, scale); + __x = u.fp8[0]; +#endif + } + + explicit __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp8_e5m2(const __amd_fp16_storage_t in, + const unsigned int seed, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + union u { + unsigned int ui32; + __amd_fp8_storage_t fp8[4]; + } u{0}; + u.ui32 = + __builtin_amdgcn_cvt_scalef32_sr_bf8_f16(u.ui32, in, seed, __amd_scale_to_float(scale), 0); + __x = u.fp8[0]; +#else + using namespace fcbx; + union { + uint32_t ui32; + __amd_fp8_storage_t fp8[4]; + } u{0}; + u.ui32 = from_float_sr<__amd_fp16_storage_t, Encoding::E5M2, true>(in, seed, scale); + __x = u.fp8[0]; +#endif + } + + explicit __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp8_e5m2(const __amd_bf16_storage_t in, + const unsigned int seed, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + union u { + unsigned int ui32; + __amd_fp8_storage_t fp8[4]; + } u{0}; + u.ui32 = + __builtin_amdgcn_cvt_scalef32_sr_bf8_bf16(u.ui32, in, seed, __amd_scale_to_float(scale), 0); + __x = u.fp8[0]; +#else + using namespace fcbx; + union { + uint32_t ui32; + __amd_fp8_storage_t fp8[4]; + } u{0}; + u.ui32 = from_float_sr<__amd_bf16_storage_t, Encoding::E5M2, true>(in, seed, scale); + __x = u.fp8[0]; +#endif + } + + __OCP_FP_HOST_DEVICE__ __amd_fp16_storage_t get_scaled_fp16(const __amd_scale_t scale) const { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + __amd_fp16x2_storage_t ret; + ret = __builtin_amdgcn_cvt_scalef32_f16_bf8(ret, __x, __amd_scale_to_float(scale), 0, false); + return ret[0]; +#else + using namespace fcbx; + return to_float<__amd_fp16_storage_t, Encoding::E5M2, true>(__x, scale); +#endif + } + + __OCP_FP_HOST_DEVICE__ __amd_bf16_storage_t get_scaled_bf16(const __amd_scale_t scale) const { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + static_assert(sizeof(unsigned int) == sizeof(__amd_fp8_storage_t[4])); + union { + __amd_fp8_storage_t fp8[4]; + unsigned int ui32; + } u; + u.fp8[0] = __x; + auto ret = + __builtin_amdgcn_cvt_scalef32_pk_bf16_fp8(u.ui32, __amd_scale_to_float(scale), false); + return ret[0]; +#else + using namespace fcbx; + return to_float<__amd_bf16_storage_t, Encoding::E5M2, true>(__x, scale); +#endif + } + + __OCP_FP_HOST_DEVICE__ float get_scaled_float(const __amd_scale_t scale) const { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + return __builtin_amdgcn_cvt_scalef32_f32_bf8(__x, __amd_scale_to_float(scale), 0); +#else + using namespace fcbx; + return to_float(__x, scale); +#endif + } + + __OCP_FP_HOST_DEVICE__ operator float() const { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + return __builtin_amdgcn_cvt_f32_bf8(__x, 0); +#else + using namespace fcbx; + return to_float(__x, 0); +#endif + } +}; + +struct __hipext_ocp_fp8x2_e4m3 { + __amd_fp8x2_storage_t __x; + static const __amd_fp8_interpretation_t __default_interpret = __AMD_OCP_E4M3; + + __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp8x2_e4m3() = default; + + explicit __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp8x2_e4m3(const float a, const float b) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + union { + int i32; + __amd_fp8x2_storage_t fp8x2[2]; + } u{0}; + u.i32 = __builtin_amdgcn_cvt_pk_fp8_f32(a, b, u.i32, false); + __x = u.fp8x2[0]; +#else + using namespace fcbx; + union { + uint32_t ui32; + __amd_fp8x2_storage_t fp8x2[2]; + } u{0}; + u.ui32 = from_float(b, 0 /*scale*/); + u.ui32 <<= 8; + u.ui32 |= from_float(a, 0 /*scale*/); + __x = u.fp8x2[0]; +#endif + } + + explicit __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp8x2_e4m3(const __amd_floatx2_storage_t in) + : __hipext_ocp_fp8x2_e4m3(in[0], in[1]) {} + + explicit __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp8x2_e4m3(const float a, const float b, + __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + union { + __amd_shortx2_storage_t shortx2; + __amd_fp8x2_storage_t fp8x2[2]; + } u{0}; + u.shortx2 = __builtin_amdgcn_cvt_scalef32_pk_fp8_f32(u.shortx2, a, b, + __amd_scale_to_float(scale), false); + __x = u.fp8x2[0]; +#else + using namespace fcbx; + union { + uint32_t ui32; + __amd_fp8x2_storage_t fp8x2[2]; + } u{0}; + u.ui32 = from_float(b, scale); + u.ui32 <<= 8; + u.ui32 |= from_float(a, scale); + __x = u.fp8x2[0]; +#endif + } + + explicit __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp8x2_e4m3(const __amd_floatx2_storage_t in, + const __amd_scale_t scale) + : __hipext_ocp_fp8x2_e4m3(in[0], in[1], scale) {} + + + __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp8x2_e4m3(const __amd_fp16x2_storage_t in, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + union { + __amd_shortx2_storage_t shortx2; + __amd_fp8x2_storage_t fp8x2[2]; + } u{0}; + u.shortx2 = + __builtin_amdgcn_cvt_scalef32_pk_fp8_f16(u.shortx2, in, __amd_scale_to_float(scale), false); + __x = u.fp8x2[0]; +#else + using namespace fcbx; + union { + uint32_t ui32; + __amd_fp8x2_storage_t fp8x2[2]; + } u{0}; + u.ui32 = from_float<__amd_fp16_storage_t, Encoding::E4M3, true>(in[1], scale); + u.ui32 <<= 8; + u.ui32 |= from_float<__amd_fp16_storage_t, Encoding::E4M3, true>(in[0], scale); + __x = u.fp8x2[0]; +#endif + } + + explicit __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp8x2_e4m3(const __amd_bf16x2_storage_t in, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + union { + __amd_shortx2_storage_t shortx2; + __amd_fp8x2_storage_t fp8x2[2]; + } u{0}; + u.shortx2 = __builtin_amdgcn_cvt_scalef32_pk_fp8_bf16(u.shortx2, in, + __amd_scale_to_float(scale), false); + __x = u.fp8x2[0]; +#else + using namespace fcbx; + union { + uint32_t ui32; + __amd_fp8x2_storage_t fp8x2[2]; + } u{0}; + u.ui32 = from_float<__amd_bf16_storage_t, Encoding::E4M3, true>(in[1], scale); + u.ui32 <<= 8; + u.ui32 |= from_float<__amd_bf16_storage_t, Encoding::E4M3, true>(in[0], scale); + __x = u.fp8x2[0]; +#endif + } + + __OCP_FP_HOST_DEVICE__ __amd_fp16x2_storage_t get_scaled_fp16x2(const __amd_scale_t scale) const { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + union { + __amd_fp8x2_storage_t fp8x2[2]; + unsigned int ui32; + } u; + u.fp8x2[0] = __x; + return __builtin_amdgcn_cvt_scalef32_pk_f16_fp8(u.ui32, __amd_scale_to_float(scale), false); +#else + using namespace fcbx; + __amd_fp16x2_storage_t ret; + ret[0] = to_float<__amd_fp16_storage_t, Encoding::E4M3, true>(__x & 0xFF, scale); + ret[1] = to_float<__amd_fp16_storage_t, Encoding::E4M3, true>(__x >> 8, scale); + return ret; +#endif + } + + __OCP_FP_HOST_DEVICE__ __amd_bf16x2_storage_t get_scaled_bf16x2(const __amd_scale_t scale) const { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + union { + __amd_fp8x2_storage_t fp8x2[2]; + unsigned int ui32; + } u; + u.fp8x2[0] = __x; + return __builtin_amdgcn_cvt_scalef32_pk_bf16_fp8(u.ui32, __amd_scale_to_float(scale), false); +#else + using namespace fcbx; + __amd_bf16x2_storage_t ret; + ret[0] = to_float<__amd_bf16_storage_t, Encoding::E4M3, true>(__x & 0xFF, scale); + ret[1] = to_float<__amd_bf16_storage_t, Encoding::E4M3, true>(__x >> 8, scale); + return ret; +#endif + } + + __OCP_FP_HOST_DEVICE__ __amd_floatx2_storage_t + get_scaled_floatx2(const __amd_scale_t scale) const { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + return __builtin_amdgcn_cvt_scalef32_pk_f32_fp8(__x, __amd_scale_to_float(scale), false); +#else + using namespace fcbx; + __amd_floatx2_storage_t ret; + ret[0] = to_float(__x & 0xFF, scale); + ret[1] = to_float(__x >> 8, scale); + return ret; +#endif + } + + __OCP_FP_HOST_DEVICE__ + operator __amd_floatx2_storage_t() const { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + return __builtin_amdgcn_cvt_pk_f32_fp8(__x, false); +#else + using namespace fcbx; + __amd_floatx2_storage_t ret; + ret[0] = to_float(__x & 0xFF, 0); + ret[1] = to_float(__x >> 8, 0); + return ret; +#endif + } +}; + +struct __hipext_ocp_fp8x2_e5m2 { + __amd_fp8x2_storage_t __x; + static const __amd_fp8_interpretation_t __default_interpret = __AMD_OCP_E5M2; + + __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp8x2_e5m2() = default; + + __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp8x2_e5m2(const float a, const float b) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + union { + int i32; + __amd_fp8x2_storage_t fp8x2[2]; + } u{0}; + u.i32 = __builtin_amdgcn_cvt_pk_bf8_f32(a, b, u.i32, false); + __x = u.fp8x2[0]; +#else + using namespace fcbx; + union { + uint32_t ui32; + __amd_fp8x2_storage_t fp8x2[2]; + } u{0}; + u.ui32 = from_float(b, 0 /*scale*/); + u.ui32 <<= 8; + u.ui32 |= from_float(a, 0 /*scale*/); + __x = u.fp8x2[0]; +#endif + } + + __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp8x2_e5m2(const __amd_floatx2_storage_t in) + : __hipext_ocp_fp8x2_e5m2(in[0], in[1]) {} + + __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp8x2_e5m2(const float a, const float b, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + union { + __amd_shortx2_storage_t shortx2; + __amd_fp8x2_storage_t fp8x2[2]; + } u{0}; + u.shortx2 = __builtin_amdgcn_cvt_scalef32_pk_bf8_f32(u.shortx2, a, b, + __amd_scale_to_float(scale), false); + __x = u.fp8x2[0]; +#else + using namespace fcbx; + union { + uint32_t ui32; + __amd_fp8x2_storage_t fp8x2[2]; + } u{0}; + u.ui32 = from_float(b, scale); + u.ui32 <<= 8; + u.ui32 |= from_float(a, scale); + __x = u.fp8x2[0]; +#endif + } + + __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp8x2_e5m2(const __amd_floatx2_storage_t in, + const __amd_scale_t scale) + : __hipext_ocp_fp8x2_e5m2(in[0], in[1], scale) {} + + __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp8x2_e5m2(const __amd_fp16x2_storage_t in, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + union { + __amd_shortx2_storage_t shortx2; + __amd_fp8x2_storage_t fp8x2[2]; + } u{0}; + u.shortx2 = + __builtin_amdgcn_cvt_scalef32_pk_bf8_f16(u.shortx2, in, __amd_scale_to_float(scale), false); + __x = u.fp8x2[0]; +#else + using namespace fcbx; + union { + uint32_t ui32; + __amd_fp8x2_storage_t fp8x2[2]; + } u{0}; + u.ui32 = from_float<__amd_fp16_storage_t, Encoding::E5M2, true>(in[1], scale); + u.ui32 <<= 8; + u.ui32 |= from_float<__amd_fp16_storage_t, Encoding::E5M2, true>(in[0], scale); + __x = u.fp8x2[0]; +#endif + } + + __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp8x2_e5m2(const __amd_bf16x2_storage_t in, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + union { + __amd_shortx2_storage_t shortx2; + __amd_fp8x2_storage_t fp8x2[2]; + } u{0}; + u.shortx2 = __builtin_amdgcn_cvt_scalef32_pk_bf8_bf16(u.shortx2, in, + __amd_scale_to_float(scale), false); + __x = u.fp8x2[0]; +#else + using namespace fcbx; + union { + uint32_t ui32; + __amd_fp8x2_storage_t fp8x2[2]; + } u{0}; + u.ui32 = from_float<__amd_bf16_storage_t, Encoding::E5M2, true>(in[1], scale); + u.ui32 <<= 8; + u.ui32 |= from_float<__amd_bf16_storage_t, Encoding::E5M2, true>(in[0], scale); + __x = u.fp8x2[0]; +#endif + } + + __OCP_FP_HOST_DEVICE__ __amd_fp16x2_storage_t get_scaled_fp16x2(const __amd_scale_t scale) const { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + union { + __amd_fp8x2_storage_t fp8x2[2]; + unsigned int ui32; + } u; + u.fp8x2[0] = __x; + return __builtin_amdgcn_cvt_scalef32_pk_f16_bf8(u.ui32, __amd_scale_to_float(scale), false); +#else + using namespace fcbx; + __amd_fp16x2_storage_t ret; + ret[0] = to_float<__amd_fp16_storage_t, Encoding::E5M2, true>(__x & 0xFF, scale); + ret[1] = to_float<__amd_fp16_storage_t, Encoding::E5M2, true>(__x >> 8, scale); + return ret; +#endif + } + + __OCP_FP_HOST_DEVICE__ __amd_bf16x2_storage_t get_scaled_bf16x2(const __amd_scale_t scale) const { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + union { + __amd_fp8x2_storage_t fp8x2[2]; + unsigned int ui32; + } u; + u.fp8x2[0] = __x; + return __builtin_amdgcn_cvt_scalef32_pk_bf16_bf8(u.ui32, __amd_scale_to_float(scale), false); +#else + using namespace fcbx; + __amd_bf16x2_storage_t ret; + ret[0] = to_float<__amd_bf16_storage_t, Encoding::E5M2, true>(__x & 0xFF, scale); + ret[1] = to_float<__amd_bf16_storage_t, Encoding::E5M2, true>(__x >> 8, scale); + return ret; +#endif + } + + __OCP_FP_HOST_DEVICE__ __amd_floatx2_storage_t + get_scaled_floatx2(const __amd_scale_t scale) const { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + return __builtin_amdgcn_cvt_scalef32_pk_f32_bf8(__x, __amd_scale_to_float(scale), false); +#else + using namespace fcbx; + __amd_floatx2_storage_t ret; + ret[0] = to_float(__x & 0xFF, scale); + ret[1] = to_float(__x >> 8, scale); + return ret; +#endif + } + + __OCP_FP_HOST_DEVICE__ + operator __amd_floatx2_storage_t() const { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + return __builtin_amdgcn_cvt_pk_f32_bf8(__x, false); +#else + using namespace fcbx; + __amd_floatx2_storage_t ret; + ret[0] = to_float(__x & 0xFF, 0); + ret[1] = to_float(__x >> 8, 0); + return ret; +#endif + } +}; + +struct __hipext_ocp_fp6x32_e2m3 { + __amd_fp6x32_storage_t __x; + static const __amd_fp6_interpretation_t __default_interpret = __AMD_OCP_E2M3; + + __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp6x32_e2m3(const __amd_floatx16_storage_t in1, + const __amd_floatx16_storage_t in2, + const __amd_scale_t scale) +#if HIP_ENABLE_GFX950_OCP_BUILTINS + : __x(__builtin_amdgcn_cvt_scalef32_2xpk16_fp6_f32(in1, in2, __amd_scale_to_float(scale))) { + } +#else + { + using namespace fcbx; + __amd_floatx32_storage_t tmp_in; + for (size_t i = 0; i < 16; i++) { + tmp_in[i] = in1[i]; + tmp_in[16 + i] = in2[i]; + } + __x = fp6_cvt_packedx32<__amd_floatx32_storage_t, __amd_fp6x32_storage_t, float, + Encoding::IEEE754, Encoding::E2M3>(tmp_in, scale); + } +#endif + + __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp6x32_e2m3(const __amd_floatx32_storage_t in, + const unsigned int round, + const __amd_scale_t scale) +#if HIP_ENABLE_GFX950_OCP_BUILTINS + : __x(__builtin_amdgcn_cvt_scalef32_sr_pk32_fp6_f32(in, round, __amd_scale_to_float(scale))){} +#else + : __x(fcbx::fp6_cvt_packedx32<__amd_floatx32_storage_t, __amd_fp6x32_storage_t, float, + fcbx::Encoding::IEEE754, fcbx::Encoding::E2M3>(in, scale)) { + } +#endif + + __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp6x32_e2m3( + const __amd_fp16x32_storage_t in, const unsigned int round, const __amd_scale_t scale) +#if HIP_ENABLE_GFX950_OCP_BUILTINS + : __x(__builtin_amdgcn_cvt_scalef32_sr_pk32_fp6_f16(in, round, __amd_scale_to_float(scale))){} +#else + : __x(fcbx::fp6_cvt_packedx32<__amd_fp16x32_storage_t, __amd_fp6x32_storage_t, + __amd_fp16_storage_t, fcbx::Encoding::E5M10, + fcbx::Encoding::E2M3, true>(in, scale, round)) { + } +#endif + + __OCP_FP_HOST_DEVICE__ + __hipext_ocp_fp6x32_e2m3(const __amd_fp16x32_storage_t in, const __amd_scale_t scale) +#if HIP_ENABLE_GFX950_OCP_BUILTINS + : __x(__builtin_amdgcn_cvt_scalef32_pk32_fp6_f16(in, __amd_scale_to_float(scale))){} +#else + : __x(fcbx::fp6_cvt_packedx32<__amd_fp16x32_storage_t, __amd_fp6x32_storage_t, + __amd_fp16_storage_t, fcbx::Encoding::E5M10, + fcbx::Encoding::E2M3>(in, scale)) { + } +#endif + + __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp6x32_e2m3( + const __amd_bf16x32_storage_t in, const unsigned int round, const __amd_scale_t scale) +#if HIP_ENABLE_GFX950_OCP_BUILTINS + : __x(__builtin_amdgcn_cvt_scalef32_sr_pk32_fp6_bf16(in, round, + __amd_scale_to_float(scale))){} +#else + : __x(fcbx::fp6_cvt_packedx32<__amd_bf16x32_storage_t, __amd_fp6x32_storage_t, + __amd_bf16_storage_t, fcbx::Encoding::E8M7, + fcbx::Encoding::E2M3, true>(in, scale, round)) { + } +#endif + + __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp6x32_e2m3(const __amd_bf16x32_storage_t in, + const __amd_scale_t scale) +#if HIP_ENABLE_GFX950_OCP_BUILTINS + : __x(__builtin_amdgcn_cvt_scalef32_pk32_fp6_bf16(in, __amd_scale_to_float(scale))){} +#else + : __x(fcbx::fp6_cvt_packedx32<__amd_bf16x32_storage_t, __amd_fp6x32_storage_t, + __amd_bf16_storage_t, fcbx::Encoding::E8M7, + fcbx::Encoding::E2M3>(in, scale)) { + } +#endif + + __OCP_FP_HOST_DEVICE__ __amd_floatx32_storage_t + get_scaled_floatx32(const __amd_scale_t scale) const { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + return __builtin_amdgcn_cvt_scalef32_pk32_f32_fp6(__x, __amd_scale_to_float(scale)); +#else + using namespace fcbx; + return fp6_cvt_packedx32<__amd_fp6x32_storage_t, __amd_floatx32_storage_t, float, + Encoding::E2M3, Encoding::IEEE754>(__x, scale); +#endif + } + + __OCP_FP_HOST_DEVICE__ __amd_fp16x32_storage_t + get_scaled_fp16x32(const __amd_scale_t scale) const { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + return __builtin_amdgcn_cvt_scalef32_pk32_f16_fp6(__x, __amd_scale_to_float(scale)); +#else + using namespace fcbx; + return fp6_cvt_packedx32<__amd_fp6x32_storage_t, __amd_fp16x32_storage_t, __amd_fp16_storage_t, + Encoding::E2M3, Encoding::E5M10>(__x, scale); +#endif + } + + __OCP_FP_HOST_DEVICE__ __amd_bf16x32_storage_t + get_scaled_bf16x32(const __amd_scale_t scale) const { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + return __builtin_amdgcn_cvt_scalef32_pk32_bf16_fp6(__x, __amd_scale_to_float(scale)); +#else + using namespace fcbx; + return fp6_cvt_packedx32<__amd_fp6x32_storage_t, __amd_bf16x32_storage_t, __amd_bf16_storage_t, + Encoding::E2M3, Encoding::E8M7>(__x, scale); +#endif + } +}; + +struct __hipext_ocp_fp6x32_e3m2 { + __amd_fp6x32_storage_t __x; + static const __amd_fp6_interpretation_t __default_interpret = __AMD_OCP_E3M2; + + __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp6x32_e3m2(const __amd_floatx16_storage_t in1, + const __amd_floatx16_storage_t in2, + const __amd_scale_t scale) +#if HIP_ENABLE_GFX950_OCP_BUILTINS + : __x(__builtin_amdgcn_cvt_scalef32_2xpk16_bf6_f32(in1, in2, __amd_scale_to_float(scale))) { + } +#else + { + using namespace fcbx; + __amd_floatx32_storage_t tmp_in; + for (size_t i = 0; i < 16; i++) { + tmp_in[i] = in1[i]; + tmp_in[16 + i] = in2[i]; + } + __x = fp6_cvt_packedx32<__amd_floatx32_storage_t, __amd_fp6x32_storage_t, float, + Encoding::IEEE754, Encoding::E3M2>(tmp_in, scale); + } +#endif + + __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp6x32_e3m2(const __amd_floatx32_storage_t in, + const unsigned int round, + const __amd_scale_t scale) +#if HIP_ENABLE_GFX950_OCP_BUILTINS + : __x(__builtin_amdgcn_cvt_scalef32_sr_pk32_bf6_f32(in, round, __amd_scale_to_float(scale))){} +#else + : __x(fcbx::fp6_cvt_packedx32<__amd_floatx32_storage_t, __amd_fp6x32_storage_t, float, + fcbx::Encoding::IEEE754, fcbx::Encoding::E3M2>(in, scale)) { + } +#endif + + __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp6x32_e3m2( + const __amd_fp16x32_storage_t in, const unsigned int round, const __amd_scale_t scale) +#if HIP_ENABLE_GFX950_OCP_BUILTINS + : __x(__builtin_amdgcn_cvt_scalef32_sr_pk32_bf6_f16(in, round, __amd_scale_to_float(scale))){} +#else + : __x(fcbx::fp6_cvt_packedx32<__amd_fp16x32_storage_t, __amd_fp6x32_storage_t, + __amd_fp16_storage_t, fcbx::Encoding::E5M10, + fcbx::Encoding::E3M2, true>(in, scale, round)) { + } +#endif + + __OCP_FP_HOST_DEVICE__ + __hipext_ocp_fp6x32_e3m2(const __amd_fp16x32_storage_t in, const __amd_scale_t scale) +#if HIP_ENABLE_GFX950_OCP_BUILTINS + : __x(__builtin_amdgcn_cvt_scalef32_pk32_bf6_f16(in, __amd_scale_to_float(scale))){} +#else + : __x(fcbx::fp6_cvt_packedx32<__amd_fp16x32_storage_t, __amd_fp6x32_storage_t, + __amd_fp16_storage_t, fcbx::Encoding::E5M10, + fcbx::Encoding::E3M2>(in, scale)) { + } +#endif + + __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp6x32_e3m2( + const __amd_bf16x32_storage_t in, const unsigned int round, const __amd_scale_t scale) +#if HIP_ENABLE_GFX950_OCP_BUILTINS + : __x(__builtin_amdgcn_cvt_scalef32_sr_pk32_bf6_bf16(in, round, + __amd_scale_to_float(scale))){} +#else + : __x(fcbx::fp6_cvt_packedx32<__amd_bf16x32_storage_t, __amd_fp6x32_storage_t, + __amd_bf16_storage_t, fcbx::Encoding::E8M7, + fcbx::Encoding::E3M2, true>(in, scale, round)) { + } +#endif + + __OCP_FP_HOST_DEVICE__ + __hipext_ocp_fp6x32_e3m2(const __amd_bf16x32_storage_t in, const __amd_scale_t scale) +#if HIP_ENABLE_GFX950_OCP_BUILTINS + : __x(__builtin_amdgcn_cvt_scalef32_pk32_bf6_bf16(in, __amd_scale_to_float(scale))){} +#else + : __x(fcbx::fp6_cvt_packedx32<__amd_bf16x32_storage_t, __amd_fp6x32_storage_t, + __amd_bf16_storage_t, fcbx::Encoding::E8M7, + fcbx::Encoding::E3M2>(in, scale)) { + } +#endif + + __OCP_FP_HOST_DEVICE__ __amd_floatx32_storage_t + get_scaled_floatx32(const __amd_scale_t scale) const { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + return __builtin_amdgcn_cvt_scalef32_pk32_f32_bf6(__x, __amd_scale_to_float(scale)); +#else + using namespace fcbx; + return fp6_cvt_packedx32<__amd_fp6x32_storage_t, __amd_floatx32_storage_t, float, + Encoding::E3M2, Encoding::IEEE754>(__x, scale); +#endif + } + + __OCP_FP_HOST_DEVICE__ __amd_fp16x32_storage_t + get_scaled_fp16x32(const __amd_scale_t scale) const { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + return __builtin_amdgcn_cvt_scalef32_pk32_f16_bf6(__x, __amd_scale_to_float(scale)); +#else + using namespace fcbx; + return fp6_cvt_packedx32<__amd_fp6x32_storage_t, __amd_fp16x32_storage_t, __amd_fp16_storage_t, + Encoding::E3M2, Encoding::E5M10>(__x, scale); +#endif + } + + __OCP_FP_HOST_DEVICE__ __amd_bf16x32_storage_t + get_scaled_bf16x32(const __amd_scale_t scale) const { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + return __builtin_amdgcn_cvt_scalef32_pk32_bf16_bf6(__x, __amd_scale_to_float(scale)); +#else + using namespace fcbx; + return fp6_cvt_packedx32<__amd_fp6x32_storage_t, __amd_bf16x32_storage_t, __amd_bf16_storage_t, + Encoding::E3M2, Encoding::E8M7>(__x, scale); +#endif + } +}; + +struct __hipext_ocp_fp4x2_e2m1 { + __amd_fp4x2_storage_t __x; + static const __amd_fp4_interpretation_t __default_interpret = __AMD_OCP_E2M1; + + static_assert(sizeof(unsigned int) == sizeof(__amd_fp4x2_storage_t[4])); + + __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp4x2_e2m1(const float a, const float b, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + union { + unsigned int ui32; + __amd_fp4x2_storage_t fp4x2[4]; + } u{0}; + u.ui32 = __builtin_amdgcn_cvt_scalef32_pk_fp4_f32(u.ui32, a, b, __amd_scale_to_float(scale), 0); + __x = u.fp4x2[0]; +#else + using namespace fcbx; + auto l = from_float(a, scale); + auto r = from_float(b, scale); + __x = r << 4 | l; +#endif + } + + __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp4x2_e2m1(const __amd_floatx2_storage_t in, + const __amd_scale_t scale) + : __hipext_ocp_fp4x2_e2m1(in[0], in[1], scale) {} + + __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp4x2_e2m1(const __amd_bf16x2_storage_t in, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + union { + unsigned int ui32; + __amd_fp4x2_storage_t fp4x2[4]; + } u{0}; + u.ui32 = __builtin_amdgcn_cvt_scalef32_pk_fp4_bf16(u.ui32, in, __amd_scale_to_float(scale), 1); + __x = u.fp4x2[1]; +#else + using namespace fcbx; + auto l = from_float<__amd_bf16_storage_t, Encoding::E2M1, true>(in[0], scale); + auto r = from_float<__amd_bf16_storage_t, Encoding::E2M1, true>(in[1], scale); + __x = r << 4 | l; +#endif + } + + __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp4x2_e2m1(const __amd_fp16x2_storage_t in, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + union { + unsigned int ui32; + __amd_fp4x2_storage_t fp4x2[4]; + } u{0}; + u.ui32 = __builtin_amdgcn_cvt_scalef32_pk_fp4_f16(u.ui32, in, __amd_scale_to_float(scale), 1); + __x = u.fp4x2[1]; +#else + using namespace fcbx; + auto l = from_float<__amd_fp16_storage_t, Encoding::E2M1, true>(in[0], scale); + auto r = from_float<__amd_fp16_storage_t, Encoding::E2M1, true>(in[1], scale); + __x = r << 4 | l; +#endif + } + + __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp4x2_e2m1(const __amd_floatx2_storage_t in, + const unsigned int seed, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + union { + unsigned int ui32; + __amd_fp4x2_storage_t fp4x2[4]; + } u{0}; + u.ui32 = __builtin_amdgcn_cvt_scalef32_sr_pk_fp4_f32(u.ui32, in, seed, + __amd_scale_to_float(scale), 1); + __x = u.fp4x2[1]; +#else + static_assert(sizeof(__amd_fp4x2_storage_t[4]) == sizeof(uint32_t)); + union u { + uint32_t ui32t; + __amd_fp4x2_storage_t fp4x2[4]; + } u{0}; + using namespace fcbx; + auto l = from_float_sr(in[0], seed, scale); + auto r = from_float_sr(in[1], seed, scale); + r <<= 4; + l |= r; + u.ui32t = l; + __x = u.fp4x2[0]; +#endif + } + + __OCP_FP_HOST_DEVICE__ + __hipext_ocp_fp4x2_e2m1(const __amd_bf16x2_storage_t in, const unsigned int seed, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + union { + unsigned int ui32; + __amd_fp4x2_storage_t fp4x2[4]; + } u{0}; + u.ui32 = __builtin_amdgcn_cvt_scalef32_sr_pk_fp4_bf16(u.ui32, in, seed, + __amd_scale_to_float(scale), 1); + __x = u.fp4x2[1]; +#else + static_assert(sizeof(__amd_fp4x2_storage_t[4]) == sizeof(uint32_t)); + union u { + uint32_t ui32t; + __amd_fp4x2_storage_t fp4x2[4]; + } u{0}; + using namespace fcbx; + auto l = from_float_sr<__amd_bf16_storage_t, Encoding::E2M1, true>(in[0], seed, scale); + auto r = from_float_sr<__amd_bf16_storage_t, Encoding::E2M1, true>(in[1], seed, scale); + r <<= 4; + l |= r; + u.ui32t = l; + __x = u.fp4x2[0]; +#endif + } + + __OCP_FP_HOST_DEVICE__ __hipext_ocp_fp4x2_e2m1(const __amd_fp16x2_storage_t in, + const unsigned int seed, + const __amd_scale_t scale) { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + union { + unsigned int ui32; + __amd_fp4x2_storage_t fp4x2[4]; + } u{0}; + u.ui32 = __builtin_amdgcn_cvt_scalef32_sr_pk_fp4_f16(u.ui32, in, seed, + __amd_scale_to_float(scale), 1); + __x = u.fp4x2[1]; +#else + static_assert(sizeof(__amd_fp4x2_storage_t[4]) == sizeof(uint32_t)); + union u { + uint32_t ui32t; + __amd_fp4x2_storage_t fp4x2[4]; + } u{0}; + using namespace fcbx; + auto l = from_float_sr<__amd_fp16_storage_t, Encoding::E2M1, true>(in[0], seed, scale); + auto r = from_float_sr<__amd_fp16_storage_t, Encoding::E2M1, true>(in[1], seed, scale); + r <<= 4; + l |= r; + u.ui32t = l; + __x = u.fp4x2[0]; +#endif + } + + __OCP_FP_HOST_DEVICE__ __amd_floatx2_storage_t + get_scaled_floatx2(const __amd_scale_t scale) const { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + return __builtin_amdgcn_cvt_scalef32_pk_f32_fp4(__x, __amd_scale_to_float(scale), 0); +#else + using namespace fcbx; + __amd_floatx2_storage_t ret{to_float(__x & 0xFu, scale), + to_float(__x >> 4, scale)}; + return ret; +#endif + } + + __OCP_FP_HOST_DEVICE__ __amd_fp16x2_storage_t get_scaled_fp16x2(const __amd_scale_t scale) const { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + return __builtin_amdgcn_cvt_scalef32_pk_f16_fp4(__x, __amd_scale_to_float(scale), 0); +#else + using namespace fcbx; + __amd_fp16x2_storage_t ret{ + to_float<__amd_fp16_storage_t, Encoding::E2M1, true>(__x & 0xFu, scale), + to_float<__amd_fp16_storage_t, Encoding::E2M1, true>(__x >> 4, scale)}; + return ret; +#endif + } + + __OCP_FP_HOST_DEVICE__ __amd_bf16x2_storage_t get_scaled_bf16x2(const __amd_scale_t scale) const { +#if HIP_ENABLE_GFX950_OCP_BUILTINS + return __builtin_amdgcn_cvt_scalef32_pk_bf16_fp4(__x, __amd_scale_to_float(scale), 0); +#else + using namespace fcbx; + __amd_bf16x2_storage_t ret{ + to_float<__amd_bf16_storage_t, Encoding::E2M1, true>(__x & 0xFu, scale), + to_float<__amd_bf16_storage_t, Encoding::E2M1, true>(__x >> 4, scale)}; + return ret; +#endif + } +}; diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_ocp_host.hpp b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_ocp_host.hpp new file mode 100644 index 0000000000..79710875b9 --- /dev/null +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_ocp_host.hpp @@ -0,0 +1,961 @@ +#pragma once + +#include "amd_hip_ocp_types.h" + +#if !defined(__HIPCC_RTC__) +#include +#include +#include +#include +#endif + +namespace fcbx { +constexpr int8_t OCP_SCALE_EXP_NAN = -128; + +enum class Encoding : size_t { + E2M1 = 0, + E2M3, + E3M2, + E4M3, + E4M3Mx, + E4M3Nanoo, + E5M2, + E5M2Mx, + E5M2Nanoo, + + E5M10, // FP16 + E8M7, // BF16 + + IEEE754, + + // Keep this one last + NumEncodings, +}; +enum fp16 : uint16_t {}; +enum bf16 : uint16_t {}; + +struct Float { + int32_t ExpBias; + uint32_t ExpBits; + uint32_t ExpMask; + uint32_t ManBits; + uint32_t ManMask; + int32_t MaxExp; + int32_t MinExp; + bool MxScale; + bool HasNaN; + bool HasInf; +}; + +static const float ieee754_nan = std::numeric_limits::quiet_NaN(); +static const float ieee754_inf = std::numeric_limits::infinity(); + +__OCP_FP_HOST_DEVICE_STATIC__ uint32_t U32(float f) { + static_assert(sizeof(uint32_t) == sizeof(float)); + union { + float f32; + uint32_t ui32; + } u{f}; + return u.ui32; +} + +__OCP_FP_HOST_DEVICE_STATIC__ float F32(uint32_t u32) { + static_assert(sizeof(uint32_t) == sizeof(float)); + union { + uint32_t ui32; + float f32; + } u{u32}; + return u.f32; +} + +constexpr __OCP_FP_HOST_DEVICE_STATIC__ uint32_t bitmask(uint32_t bits) { + if (bits < 1) return 0; + return ((uint32_t)1 << bits) - 1; +} + +constexpr std::array init() { + std::array a{}; + + a[(size_t)Encoding::E2M1] = { + .ExpBias = 1, + .ExpBits = 2, + .ExpMask = bitmask(2), + .ManBits = 1, + .ManMask = bitmask(1), + .MaxExp = 2, + .MinExp = 0, + .MxScale = true, + .HasNaN = false, + .HasInf = false, + }; + + a[(size_t)Encoding::E2M3] = { + .ExpBias = 1, + .ExpBits = 2, + .ExpMask = bitmask(2), + .ManBits = 3, + .ManMask = bitmask(3), + .MaxExp = 2, + .MinExp = 0, + .MxScale = true, + .HasNaN = false, + .HasInf = false, + }; + + a[(size_t)Encoding::E3M2] = { + .ExpBias = 3, + .ExpBits = 3, + .ExpMask = bitmask(3), + .ManBits = 2, + .ManMask = bitmask(2), + .MaxExp = 4, + .MinExp = -2, + .MxScale = true, + .HasNaN = false, + .HasInf = false, + }; + + a[(size_t)Encoding::E4M3] = { + .ExpBias = 7, + .ExpBits = 4, + .ExpMask = bitmask(4), + .ManBits = 3, + .ManMask = bitmask(3), + .MaxExp = 8, + .MinExp = -6, + .MxScale = false, + .HasNaN = true, + .HasInf = false, + }; + + a[(size_t)Encoding::E4M3Mx] = { + .ExpBias = 7, + .ExpBits = 4, + .ExpMask = bitmask(4), + .ManBits = 3, + .ManMask = bitmask(3), + .MaxExp = 8, + .MinExp = -6, + .MxScale = true, + .HasNaN = true, + .HasInf = false, + }; + + a[(size_t)Encoding::E4M3Nanoo] = { + .ExpBias = 8, + .ExpBits = 4, + .ExpMask = bitmask(4), + .ManBits = 3, + .ManMask = bitmask(3), + .MaxExp = 7, + .MinExp = -7, + .MxScale = false, + .HasNaN = true, + .HasInf = false, + }; + + a[(size_t)Encoding::E5M2] = { + .ExpBias = 15, + .ExpBits = 5, + .ExpMask = bitmask(5), + .ManBits = 2, + .ManMask = bitmask(2), + .MaxExp = 15, + .MinExp = -14, + .MxScale = false, + .HasNaN = true, + .HasInf = true, + }; + + a[(size_t)Encoding::E5M2Mx] = { + .ExpBias = 15, + .ExpBits = 5, + .ExpMask = bitmask(5), + .ManBits = 2, + .ManMask = bitmask(2), + .MaxExp = 15, + .MinExp = -14, + .MxScale = true, + .HasNaN = true, + .HasInf = true, + }; + + a[(size_t)Encoding::E5M2Nanoo] = { + .ExpBias = 16, + .ExpBits = 5, + .ExpMask = bitmask(5), + .ManBits = 2, + .ManMask = bitmask(2), + .MaxExp = 15, + .MinExp = -15, + .MxScale = false, + .HasNaN = true, + .HasInf = true, + }; + + a[(size_t)Encoding::E5M10] = { + .ExpBias = 15, + .ExpBits = 5, + .ExpMask = bitmask(5), + .ManBits = 10, + .ManMask = bitmask(10), + .MaxExp = 15, + .MinExp = -14, + .MxScale = false, + .HasNaN = true, + .HasInf = true, + }; + + a[(size_t)Encoding::E8M7] = { + .ExpBias = 127, + .ExpBits = 8, + .ExpMask = bitmask(8), + .ManBits = 7, + .ManMask = bitmask(7), + .MaxExp = 127, + .MinExp = -126, + .MxScale = false, + .HasNaN = true, + .HasInf = true, + }; + + a[(size_t)Encoding::IEEE754] = { + .ExpBias = 127, + .ExpBits = 8, + .ExpMask = bitmask(8), + .ManBits = 23, + .ManMask = bitmask(23), + .MaxExp = 127, + .MinExp = -126, + .MxScale = false, + .HasNaN = true, + .HasInf = true, + }; + + return a; +} + +static constexpr auto encodings = init(); + +template __OCP_FP_HOST_DEVICE_STATIC__ uint32_t exponentbits(uint32_t val) { + const auto& enc = encodings[(size_t)E]; + return (val >> enc.ManBits) & enc.ExpMask; +} + +template __OCP_FP_HOST_DEVICE_STATIC__ uint32_t mantissa(uint32_t val) { + const auto& enc = encodings[(size_t)E]; + return val & enc.ManMask; +} + +template __OCP_FP_HOST_DEVICE_STATIC__ bool issubnorm(uint32_t val) { + switch (E) { + default: + return exponentbits(val) == 0 && mantissa(val) != 0; + } + + __builtin_trap(); + // Unreachable + return false; +} + +template __OCP_FP_HOST_DEVICE_STATIC__ int32_t exponent(uint32_t val) { + const auto& enc = encodings[(size_t)E]; + auto unbiased_exp = exponentbits(val); + unbiased_exp = issubnorm(val) ? 1 : unbiased_exp; + return (int32_t)unbiased_exp - enc.ExpBias; +} + +template __OCP_FP_HOST_DEVICE_STATIC__ uint32_t signbit(uint32_t val) { + const auto& enc = encodings[(size_t)E]; + return (val >> (enc.ExpBits + enc.ManBits)) & 1; +} + +template __OCP_FP_HOST_DEVICE_STATIC__ uint32_t nan(uint32_t sign) { + const auto& enc = encodings[(size_t)E]; + + switch (E) { + case Encoding::E2M1: + return (sign << (enc.ExpBits + enc.ManBits)) | 0b0111; + case Encoding::E2M3: + case Encoding::E3M2: + return (sign << (enc.ExpBits + enc.ManBits)) | 0b011111; + case Encoding::E4M3: + case Encoding::E4M3Mx: + return (sign << (enc.ExpBits + enc.ManBits)) | 0x7f; + case Encoding::E5M2: + case Encoding::E5M2Mx: + return (sign << (enc.ExpBits + enc.ManBits)) | 0x7e; + case Encoding::E4M3Nanoo: + case Encoding::E5M2Nanoo: + return 0b10000000; + case Encoding::E5M10: + case Encoding::E8M7: + return (sign << (enc.ExpBits + enc.ManBits)) | (enc.ExpMask << enc.ManBits) | enc.ManMask; + case Encoding::IEEE754: + return U32(sign ? std::copysign(ieee754_nan, -1.0F) : ieee754_nan); + default: + __builtin_trap(); + return 0; + } +} + +template __OCP_FP_HOST_DEVICE_STATIC__ uint32_t zero(uint32_t sign) { + const auto& enc = encodings[(size_t)E]; + + switch (E) { + case Encoding::E2M1: + case Encoding::E2M3: + case Encoding::E3M2: + case Encoding::E4M3: + case Encoding::E4M3Mx: + case Encoding::E5M2: + case Encoding::E5M2Mx: + case Encoding::E5M10: + case Encoding::E8M7: + return (sign << (enc.ExpBits + enc.ManBits)) | 0; + case Encoding::E4M3Nanoo: + case Encoding::E5M2Nanoo: + return 0; + case Encoding::IEEE754: + return U32(sign ? std::copysign(0.0F, -1.0F) : 0.0F); + default: + __builtin_trap(); + return 0; + } +} + +template __OCP_FP_HOST_DEVICE_STATIC__ uint32_t inf(uint32_t sign) { + const auto& enc = encodings[(size_t)E]; + + switch (E) { + case Encoding::E2M1: + case Encoding::E2M3: + case Encoding::E3M2: + return nan(sign); + case Encoding::E4M3: + case Encoding::E4M3Mx: + case Encoding::E4M3Nanoo: + case Encoding::E5M2Nanoo: + if constexpr (sat) { + sign <<= enc.ExpBits + enc.ManBits; + return sign | 0b01111111; + } + + return nan(sign); + case Encoding::E5M2: + case Encoding::E5M2Mx: + sign <<= enc.ExpBits + enc.ManBits; + if constexpr (sat) { + return sign | 0b01111011; + } + + return sign | 0b01111100; + case Encoding::E5M10: + case Encoding::E8M7: + sign <<= enc.ExpBits + enc.ManBits; + return sign | (enc.ExpMask << enc.ManBits) | 0; + case Encoding::IEEE754: + return U32(sign ? std::copysign(ieee754_inf, -1.0F) : ieee754_inf); + default: + __builtin_trap(); + return 0; + } +} + +template __OCP_FP_HOST_DEVICE_STATIC__ bool isnan(uint32_t val) { + const auto& enc = encodings[(size_t)E]; + if (!enc.HasNaN) return false; + + if constexpr (E == Encoding::E4M3Mx || E == Encoding::E4M3 || E == Encoding::E4M3Nanoo || + E == Encoding::E5M2Nanoo) + return nan(signbit(val)) == val; + + return exponentbits(val) == enc.ExpMask && mantissa(val) != 0; +} + +template __OCP_FP_HOST_DEVICE_STATIC__ bool isinf(uint32_t val) { + const auto& enc = encodings[(size_t)E]; + if (!enc.HasInf) return false; + + if constexpr (E == Encoding::E5M10 || E == Encoding::E8M7) { + return exponentbits(val) == enc.ExpMask && mantissa(val) == 0; + } + + return inf(signbit(val)) == val; +} + +template __OCP_FP_HOST_DEVICE_STATIC__ bool iszero(uint32_t val) { + return zero(signbit(val)) == val; +} + +template __OCP_FP_HOST_DEVICE_STATIC__ bool inrange(uint32_t val) { + return !(isnan(val) || isinf(val)); +} + +template __OCP_FP_HOST_DEVICE_STATIC__ T makenan(Encoding E, uint32_t sign) { + switch (E) { + case Encoding::E5M10: + return (T)nan(sign); + case Encoding::E8M7: + return (T)nan(sign); + case Encoding::IEEE754: + return (T)F32(nan(sign)); + default: + __builtin_trap(); + // Unreachable + return T(); + } +} + +template __OCP_FP_HOST_DEVICE_STATIC__ T makeinf(Encoding E, uint32_t sign) { + switch (E) { + case Encoding::E5M10: + return (T)inf(sign); + case Encoding::E8M7: + return (T)inf(sign); + case Encoding::IEEE754: + return (T)F32(inf(sign)); + default: + __builtin_trap(); + // Unreachable + return T(); + } +} + +template __OCP_FP_HOST_DEVICE_STATIC__ T makezero(Encoding E, uint32_t sign) { + switch (E) { + case Encoding::E5M10: + return (T)zero(sign); + case Encoding::E8M7: + return (T)zero(sign); + case Encoding::IEEE754: + return (T)F32(zero(sign)); + default: + __builtin_trap(); + // Unreachable + return T(); + } +} + +template +__OCP_FP_HOST_DEVICE_STATIC__ T to_float(uint32_t u32, int8_t scale_exp) { + // We do not support bf16/fp16 <-> float + static_assert(E != Encoding::IEEE754 && E != Encoding::E5M10 && E != Encoding::E8M7); + + const auto& enc = encodings[(size_t)E]; + const auto dstE = []() -> Encoding { + if constexpr (std::is_same()) + return Encoding::IEEE754; + else if constexpr (std::is_same()) + return Encoding::E5M10; + else if constexpr (std::is_same()) + return Encoding::E8M7; + else + __builtin_trap(); + }(); + const auto& dstEnc = encodings[(size_t)dstE]; + + if (isnan(u32) || (enc.MxScale && scale_exp == OCP_SCALE_EXP_NAN)) + return makenan(dstE, signbit(u32)); + + if (isinf(u32)) return makeinf(dstE, signbit(u32)); + + if (iszero(u32)) return makezero(dstE, signbit(u32)); + + auto dstMan = mantissa(u32) << (dstEnc.ManBits - enc.ManBits); + auto dstExp = (uint32_t)(exponent(u32) + dstEnc.ExpBias); + dstExp &= dstEnc.ExpMask; + + if (issubnorm(u32)) { + auto leadbit = (uint32_t)1 << dstEnc.ManBits; + while ((dstMan & leadbit) == 0) { + dstMan <<= 1; + dstExp -= 1; + } + + dstMan &= dstEnc.ManMask; + } + + auto sign = signbit(u32) << (dstEnc.ExpBits + dstEnc.ManBits); + + if (enc.MxScale) { + int32_t exp = dstExp - dstEnc.ExpBias; + int32_t tmp = exp + (int32_t)scale_exp; + size_t diff = abs(tmp - dstEnc.MinExp); + + + if (tmp < dstEnc.MinExp) { + if (diff > dstEnc.ManBits + 1) return makezero(dstE, signbit(u32)); + + dstExp = 0; // Subnormal + dstMan |= (uint32_t)1 << dstEnc.ManBits; + + auto roundBitShift = diff - 1; + auto roundBit = (dstMan & ((uint32_t)1 << roundBitShift)) != 0; + auto stickyMask = ((uint32_t)1 << roundBitShift) - 1; + auto stickyBits = dstMan & stickyMask; + auto odd = (dstMan & ((uint32_t)1 << diff)) != 0; + + dstMan >>= diff; + + if ((roundBit && stickyBits != 0) || (roundBit && odd)) { + ++dstMan; + if ((dstMan & ((uint32_t)1 << dstEnc.ManBits)) != 0) ++dstExp; + } + + dstMan &= dstEnc.ManMask; + } else { + dstExp = (uint32_t)(exp + scale_exp + dstEnc.ExpBias); + + // Overflow: return infinity (gfx950 HW behavior) + if (dstExp >= dstEnc.ExpMask) return makeinf(dstE, signbit(u32)); + + dstExp &= dstEnc.ExpMask; + } + } + + auto dst = sign | (dstExp << dstEnc.ManBits) | dstMan; + + union { + float f32; + __amd_fp16_storage_t fp16[2]; + __amd_bf16_storage_t bf16[2]; + uint32_t u32; + } u; + u.u32 = dst; + if constexpr (std::is_same()) + return u.f32; + else if constexpr (std::is_same()) + return u.fp16[0]; + else if constexpr (std::is_same()) + return u.bf16[0]; + else + __builtin_trap(); +} + +template +__OCP_FP_HOST_DEVICE_STATIC__ uint32_t from_float_sr(T f, uint32_t seed, int8_t scale_exp) { + // We do not support bf16/fp16 <-> float + static_assert(E != Encoding::IEEE754 && E != Encoding::E5M10 && E != Encoding::E8M7); + static_assert(sizeof(__amd_fp16_storage_t[2]) == sizeof(float)); + static_assert(sizeof(__amd_bf16_storage_t[2]) == sizeof(float)); + union { + float f32; + __amd_fp16_storage_t fp16[2]; + __amd_bf16_storage_t bf16[2]; + uint32_t u32; + } u; + + if constexpr (std::is_same()) + u.f32 = f; + else if constexpr (std::is_same()) + u.fp16[0] = f; + else if constexpr (std::is_same()) + u.bf16[0] = f; + else + __builtin_trap(); + + const auto& enc = encodings[(size_t)E]; + const auto srcE = []() -> Encoding { + if constexpr (std::is_same()) + return Encoding::IEEE754; + else if constexpr (std::is_same()) + return Encoding::E5M10; + else if constexpr (std::is_same()) + return Encoding::E8M7; + else + __builtin_trap(); + }(); + const auto& srcEnc = encodings[(size_t)srcE]; + + auto srcU32 = u.u32;// (srcE == Encoding::IEEE754) ? U32(f) : (uint32_t)f; + auto signBit = signbit(srcU32); + auto sign = signBit << (enc.ExpBits + enc.ManBits); + + if (isnan(srcU32) || (enc.MxScale && scale_exp == OCP_SCALE_EXP_NAN)) + return nan(signBit); + + if (isinf(srcU32)) return inf(signBit); + + if (iszero(srcU32)) return zero(signBit); + + auto srcMan = mantissa(srcU32); + auto srcExp = exponent(srcU32); + if (enc.MxScale) { + if (issubnorm(srcU32)) { + auto leadbit = (uint32_t)1 << srcEnc.ManBits; + while ((srcMan & leadbit) == 0) { + srcMan <<= 1; + srcExp -= 1; + } + + srcMan &= srcEnc.ManMask; + } + + srcExp -= scale_exp; + } + + auto exp = srcExp; + auto man = srcMan; + bool subnorm = false; + + if (exp > enc.MaxExp) { + return inf(signBit); + } else if (exp >= enc.MinExp) { + man = srcMan; + } else if (exp < enc.MinExp) { + subnorm = true; + exp = 0; + + auto diff = (uint32_t)(enc.MinExp - srcExp); + if (diff >= 32) { + man = 0; + srcMan = 0; + } else { + srcMan |= (uint32_t)1 << srcEnc.ManBits; + srcMan >>= diff; + } + + man = srcMan; + } + + // Align random value to be one past the kept mant bit + size_t sr_shift = (32 - srcEnc.ManBits) + enc.ManBits; + + // For stochastic-rounding we add the aligned random value to the + // mantissa and then truncate (RTZ). + man += seed >> sr_shift; + + // Increment exponent when mantissa overflows due to rounding + if (man >= (uint32_t)1 << srcEnc.ManBits) ++exp; + man >>= (srcEnc.ManBits - enc.ManBits); + man &= enc.ManMask; + + if (exp > enc.MaxExp) return inf(signBit); + + auto biasedExp = (uint32_t)exp; + if (!subnorm) biasedExp = (uint32_t)(exp + enc.ExpBias); + biasedExp &= enc.ExpMask; + + auto val = sign | biasedExp << enc.ManBits | man; + if (inrange(val)) + return val; + else if (man == 0 && exp == 0) + return zero(signBit); + else + return inf(signBit); +} + + +template +__OCP_FP_HOST_DEVICE_STATIC__ uint32_t from_float(T f, int8_t scale_exp) { + // We do not support bf16/fp16 <-> float + static_assert(E != Encoding::IEEE754 && E != Encoding::E5M10 && E != Encoding::E8M7); + static_assert(sizeof(__amd_fp16_storage_t[2]) == sizeof(float)); + static_assert(sizeof(__amd_bf16_storage_t[2]) == sizeof(float)); + union { + float f32; + __amd_fp16_storage_t fp16[2]; + __amd_bf16_storage_t bf16[2]; + uint32_t u32; + } u; + + if constexpr (std::is_same()) + u.f32 = f; + else if constexpr (std::is_same()) + u.fp16[0] = f; + else if constexpr (std::is_same()) + u.bf16[0] = f; + else + __builtin_trap(); + + const auto& enc = encodings[(size_t)E]; + const auto srcE = []() -> Encoding { + if constexpr (std::is_same()) + return Encoding::IEEE754; + else if constexpr (std::is_same()) + return Encoding::E5M10; + else if constexpr (std::is_same()) + return Encoding::E8M7; + else + __builtin_trap(); + }(); + const auto& srcEnc = encodings[(size_t)srcE]; + + auto srcU32 = u.u32; // (srcE == Encoding::IEEE754) ? U32(f) : (uint32_t)f; + auto signBit = signbit(srcU32); + auto sign = signBit << (enc.ExpBits + enc.ManBits); + + if (isnan(srcU32) || (enc.MxScale && scale_exp == OCP_SCALE_EXP_NAN)) + return nan(signBit); + + if (isinf(srcU32)) return inf(signBit); + + if (iszero(srcU32)) return zero(signBit); + + auto srcMan = mantissa(srcU32); + auto srcExp = exponent(srcU32); + if (enc.MxScale) { + if (issubnorm(srcU32)) { + auto leadbit = (uint32_t)1 << srcEnc.ManBits; + while ((srcMan & leadbit) == 0) { + srcMan <<= 1; + srcExp -= 1; + } + + srcMan &= srcEnc.ManMask; + } + + srcExp -= scale_exp; + } + + auto exp = srcExp; + auto man = srcMan; + uint32_t stickyBits = 0; + bool subnorm = false; + + if (exp > enc.MaxExp) { + return inf(signBit); + } else if (exp >= enc.MinExp) { + man >>= srcEnc.ManBits - enc.ManBits; + } else if (exp < enc.MinExp) { + subnorm = true; + exp = 0; + + auto diff = (uint32_t)(enc.MinExp - srcExp); + if (diff >= 32) { + man = 0; + srcMan = 0; + } else { + srcMan |= (uint32_t)1 << srcEnc.ManBits; + stickyBits = srcMan & (((uint32_t)1 << diff) - (uint32_t)1); + srcMan >>= diff; + + man = srcMan; + man >>= srcEnc.ManBits - enc.ManBits; + man &= enc.ManMask; + } + } + + auto roundBitShift = srcEnc.ManBits - (enc.ManBits + 1); + auto roundBit = ((srcMan >> roundBitShift) & 1) != 0; + stickyBits |= srcMan & (((uint32_t)1 << roundBitShift) - 1); + auto odd = (man & 1) != 0; + + if ((roundBit && stickyBits != 0) || (roundBit && odd)) { + ++man; + if ((man & ((uint32_t)1 << enc.ManBits)) != 0) ++exp; + man &= enc.ManMask; + } + + if (exp > enc.MaxExp) return inf(signBit); + + auto biasedExp = (uint32_t)exp; + if (!subnorm) biasedExp = (uint32_t)(exp + enc.ExpBias); + biasedExp &= enc.ExpMask; + + auto val = sign | biasedExp << enc.ManBits | man; + if (inrange(val)) + return val; + else if (man == 0 && exp == 0) + return zero(signBit); + else + return inf(signBit); +} + +template +__OCP_FP_HOST_DEVICE_STATIC__ OutType fp6_cvt_packedx32(InType in, int8_t scale = 0, + uint32_t seed = 0) { + // This is tightly coupled with the definitions of the amd_ocp_types + constexpr bool in_float = std::is_same::value || + std::is_same::value || + std::is_same::value; + constexpr bool out_float = std::is_same::value || + std::is_same::value || + std::is_same::value; + using other_type = std::conditional::type; + + struct fp6x32_packed { + uint8_t val1 : 6; + uint8_t val2 : 6; + uint8_t val3 : 6; + uint8_t val4 : 6; + uint8_t val5 : 6; + uint8_t val6 : 6; + uint8_t val7 : 6; + uint8_t val8 : 6; + uint8_t val9 : 6; + uint8_t val10 : 6; + uint8_t val11 : 6; + uint8_t val12 : 6; + uint8_t val13 : 6; + uint8_t val14 : 6; + uint8_t val15 : 6; + uint8_t val16 : 6; + uint8_t val17 : 6; + uint8_t val18 : 6; + uint8_t val19 : 6; + uint8_t val20 : 6; + uint8_t val21 : 6; + uint8_t val22 : 6; + uint8_t val23 : 6; + uint8_t val24 : 6; + uint8_t val25 : 6; + uint8_t val26 : 6; + uint8_t val27 : 6; + uint8_t val28 : 6; + uint8_t val29 : 6; + uint8_t val30 : 6; + uint8_t val31 : 6; + uint8_t val32 : 6; + unsigned long long padded; + } __attribute__((packed)); + + static_assert(sizeof(other_type) == sizeof(fp6x32_packed)); + union { + other_type o; + fp6x32_packed fp6; + } u; + + // TODO maybe make it simpler + if constexpr (in_float) { + if constexpr (sr) { + u.fp6.val1 = + static_cast(from_float_sr(in[0], seed, scale)); + u.fp6.val2 = + static_cast(from_float_sr(in[1], seed, scale)); + u.fp6.val3 = + static_cast(from_float_sr(in[2], seed, scale)); + u.fp6.val4 = + static_cast(from_float_sr(in[3], seed, scale)); + u.fp6.val5 = + static_cast(from_float_sr(in[4], seed, scale)); + u.fp6.val6 = + static_cast(from_float_sr(in[5], seed, scale)); + u.fp6.val7 = + static_cast(from_float_sr(in[6], seed, scale)); + u.fp6.val8 = + static_cast(from_float_sr(in[7], seed, scale)); + u.fp6.val9 = + static_cast(from_float_sr(in[8], seed, scale)); + u.fp6.val10 = + static_cast(from_float_sr(in[9], seed, scale)); + u.fp6.val11 = + static_cast(from_float_sr(in[10], seed, scale)); + u.fp6.val12 = + static_cast(from_float_sr(in[11], seed, scale)); + u.fp6.val13 = + static_cast(from_float_sr(in[12], seed, scale)); + u.fp6.val14 = + static_cast(from_float_sr(in[13], seed, scale)); + u.fp6.val15 = + static_cast(from_float_sr(in[14], seed, scale)); + u.fp6.val16 = + static_cast(from_float_sr(in[15], seed, scale)); + u.fp6.val17 = + static_cast(from_float_sr(in[16], seed, scale)); + u.fp6.val18 = + static_cast(from_float_sr(in[17], seed, scale)); + u.fp6.val19 = + static_cast(from_float_sr(in[18], seed, scale)); + u.fp6.val20 = + static_cast(from_float_sr(in[19], seed, scale)); + u.fp6.val21 = + static_cast(from_float_sr(in[20], seed, scale)); + u.fp6.val22 = + static_cast(from_float_sr(in[21], seed, scale)); + u.fp6.val23 = + static_cast(from_float_sr(in[22], seed, scale)); + u.fp6.val24 = + static_cast(from_float_sr(in[23], seed, scale)); + u.fp6.val25 = + static_cast(from_float_sr(in[24], seed, scale)); + u.fp6.val26 = + static_cast(from_float_sr(in[25], seed, scale)); + u.fp6.val27 = + static_cast(from_float_sr(in[26], seed, scale)); + u.fp6.val28 = + static_cast(from_float_sr(in[27], seed, scale)); + u.fp6.val29 = + static_cast(from_float_sr(in[28], seed, scale)); + u.fp6.val30 = + static_cast(from_float_sr(in[29], seed, scale)); + u.fp6.val31 = + static_cast(from_float_sr(in[30], seed, scale)); + u.fp6.val32 = + static_cast(from_float_sr(in[31], seed, scale)); + } else { + u.fp6.val1 = from_float(in[0], scale); + u.fp6.val2 = from_float(in[1], scale); + u.fp6.val3 = from_float(in[2], scale); + u.fp6.val4 = from_float(in[3], scale); + u.fp6.val5 = from_float(in[4], scale); + u.fp6.val6 = from_float(in[5], scale); + u.fp6.val7 = from_float(in[6], scale); + u.fp6.val8 = from_float(in[7], scale); + u.fp6.val9 = from_float(in[8], scale); + u.fp6.val10 = from_float(in[9], scale); + u.fp6.val11 = from_float(in[10], scale); + u.fp6.val12 = from_float(in[11], scale); + u.fp6.val13 = from_float(in[12], scale); + u.fp6.val14 = from_float(in[13], scale); + u.fp6.val15 = from_float(in[14], scale); + u.fp6.val16 = from_float(in[15], scale); + u.fp6.val17 = from_float(in[16], scale); + u.fp6.val18 = from_float(in[17], scale); + u.fp6.val19 = from_float(in[18], scale); + u.fp6.val20 = from_float(in[19], scale); + u.fp6.val21 = from_float(in[20], scale); + u.fp6.val22 = from_float(in[21], scale); + u.fp6.val23 = from_float(in[22], scale); + u.fp6.val24 = from_float(in[23], scale); + u.fp6.val25 = from_float(in[24], scale); + u.fp6.val26 = from_float(in[25], scale); + u.fp6.val27 = from_float(in[26], scale); + u.fp6.val28 = from_float(in[27], scale); + u.fp6.val29 = from_float(in[28], scale); + u.fp6.val30 = from_float(in[29], scale); + u.fp6.val31 = from_float(in[30], scale); + u.fp6.val32 = from_float(in[31], scale); + } + return u.o; + } else { + OutType ret; + u.o = in; + ret[0] = to_float(u.fp6.val1, scale); + ret[1] = to_float(u.fp6.val2, scale); + ret[2] = to_float(u.fp6.val3, scale); + ret[3] = to_float(u.fp6.val4, scale); + ret[4] = to_float(u.fp6.val5, scale); + ret[5] = to_float(u.fp6.val6, scale); + ret[6] = to_float(u.fp6.val7, scale); + ret[7] = to_float(u.fp6.val8, scale); + ret[8] = to_float(u.fp6.val9, scale); + ret[9] = to_float(u.fp6.val10, scale); + ret[10] = to_float(u.fp6.val11, scale); + ret[11] = to_float(u.fp6.val12, scale); + ret[12] = to_float(u.fp6.val13, scale); + ret[13] = to_float(u.fp6.val14, scale); + ret[14] = to_float(u.fp6.val15, scale); + ret[15] = to_float(u.fp6.val16, scale); + ret[16] = to_float(u.fp6.val17, scale); + ret[17] = to_float(u.fp6.val18, scale); + ret[18] = to_float(u.fp6.val19, scale); + ret[19] = to_float(u.fp6.val20, scale); + ret[20] = to_float(u.fp6.val21, scale); + ret[21] = to_float(u.fp6.val22, scale); + ret[22] = to_float(u.fp6.val23, scale); + ret[23] = to_float(u.fp6.val24, scale); + ret[24] = to_float(u.fp6.val25, scale); + ret[25] = to_float(u.fp6.val26, scale); + ret[26] = to_float(u.fp6.val27, scale); + ret[27] = to_float(u.fp6.val28, scale); + ret[28] = to_float(u.fp6.val29, scale); + ret[29] = to_float(u.fp6.val30, scale); + ret[30] = to_float(u.fp6.val31, scale); + ret[31] = to_float(u.fp6.val32, scale); + return ret; + } +} +} // namespace fcbx diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_ocp_types.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_ocp_types.h new file mode 100644 index 0000000000..26c762f50f --- /dev/null +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_ocp_types.h @@ -0,0 +1,66 @@ +#pragma once + +#define __OCP_FP_HOST__ __host__ +#define __OCP_FP_DEVICE__ __device__ +#define __OCP_FP_HOST_DEVICE__ __OCP_FP_HOST__ __OCP_FP_DEVICE__ +#define __OCP_FP_DEVICE_STATIC__ __OCP_FP_DEVICE__ static __inline__ __attribute__((always_inline)) +#define __OCP_FP_HOST_DEVICE_STATIC__ __OCP_FP_HOST_DEVICE__ static + +static_assert(sizeof(unsigned int) == 4); +static_assert(sizeof(float) == 4); +static_assert(sizeof(unsigned short) == 2); + +#if (defined(__clang__) && (__clang_major__ > 17) && defined(__HIP__)) || \ + (defined(__GNUC__) && (__GNUC__ > 13)) +static_assert(sizeof(__bf16) == 2); +static_assert(sizeof(_Float16) == 2); +#endif + +// Although we do have some abstractions of half and bfloat16, since this will be a standalone +// header which will act as a base abstraction, and will be maintained in the future, it makes sense +// to keep these vector types separate from existing implementations. We can add conversion +// functions in a different header using these functions. +typedef uint8_t __amd_fp8_storage_t; +typedef uint16_t __amd_fp8x2_storage_t; +typedef uint8_t __amd_fp4x2_storage_t; +typedef uint32_t __amd_fp4x8_storage_t; +typedef __bf16 __amd_bf16_storage_t; +typedef _Float16 __amd_fp16_storage_t; +typedef int8_t __amd_scale_t; + +#if defined(__clang__) && (__clang_major__ > 17) && defined(__HIP__) +typedef unsigned int __attribute__((ext_vector_type(2))) __amd_uintx2_storage_t; +typedef uint8_t __attribute__((ext_vector_type(8))) __amd_fp8x8_storage_t; +typedef __bf16 __attribute__((ext_vector_type(2))) __amd_bf16x2_storage_t; +typedef __bf16 __attribute__((ext_vector_type(8))) __amd_bf16x8_storage_t; +typedef __bf16 __attribute__((ext_vector_type(32))) __amd_bf16x32_storage_t; +typedef float __attribute__((ext_vector_type(2))) __amd_floatx2_storage_t; +typedef float __attribute__((ext_vector_type(8))) __amd_floatx8_storage_t; +typedef float __attribute__((ext_vector_type(16))) __amd_floatx16_storage_t; +typedef float __attribute__((ext_vector_type(32))) __amd_floatx32_storage_t; +typedef _Float16 __attribute__((ext_vector_type(2))) __amd_fp16x2_storage_t; +typedef _Float16 __attribute__((ext_vector_type(8))) __amd_fp16x8_storage_t; +typedef _Float16 __attribute__((ext_vector_type(32))) __amd_fp16x32_storage_t; +typedef uint32_t __attribute__((ext_vector_type(6))) __amd_fp6x32_storage_t; +typedef short __attribute__((ext_vector_type(2))) __amd_shortx2_storage_t; +#elif defined(__GNUC__) && (__GNUC__ > 13) +/* GCC expects vector size in bytes */ +typedef unsigned int __attribute__((vector_size(8))) __amd_uintx2_storage_t; +typedef uint8_t __attribute__((vector_size(8))) __amd_fp8x8_storage_t; +typedef __bf16 __attribute__((vector_size(4))) __amd_bf16x2_storage_t; +typedef __bf16 __attribute__((vector_size(16))) __amd_bf16x8_storage_t; +typedef __bf16 __attribute__((vector_size(64))) __amd_bf16x32_storage_t; +typedef float __attribute__((vector_size(8))) __amd_floatx2_storage_t; +typedef float __attribute__((vector_size(32))) __amd_floatx8_storage_t; +typedef float __attribute__((vector_size(64))) __amd_floatx16_storage_t; +typedef float __attribute__((vector_size(128))) __amd_floatx32_storage_t; +typedef _Float16 __attribute__((vector_size(4))) __amd_fp16x2_storage_t; +typedef _Float16 __attribute__((vector_size(16))) __amd_fp16x8_storage_t; +typedef _Float16 __attribute__((vector_size(64))) __amd_fp16x32_storage_t; +typedef uint32_t __attribute__((vector_size(24))) __amd_fp6x32_storage_t; +typedef short __attribute__((vector_size(4))) __amd_shortx2_storage_t; +#else +#error "Only supported by HIPCC or GCC >= 13." +#endif + +static_assert(sizeof(__amd_uintx2_storage_t) == sizeof(__amd_fp8x8_storage_t));