From 4ce42011b2fa14ac7c652bc82c26cb1ace386c20 Mon Sep 17 00:00:00 2001
From: mtgu0705 <mtgu@amd.com>
Date: Mon, 28 Jul 2025 03:55:32 -0500
Subject: [PATCH] added missing file

---
 include/ck_tile/core/numeric/e8m0.hpp | 86 +++++++++++++++++++++++++++
 1 file changed, 86 insertions(+)
 create mode 100644 include/ck_tile/core/numeric/e8m0.hpp

diff --git a/include/ck_tile/core/numeric/e8m0.hpp b/include/ck_tile/core/numeric/e8m0.hpp
new file mode 100644
index 0000000000..64a69dd761
--- /dev/null
+++ b/include/ck_tile/core/numeric/e8m0.hpp
@@ -0,0 +1,86 @@
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2024-2025, Advanced Micro Devices, Inc. All rights reserved.
+
+#pragma once
+
+#include "ck/utility/type.hpp"
+
+namespace ck_tile {
+
+/**
+ * @brief Unsigned representation of a conventional biased Float32 exponent.
+ *
+ * bias = 127;
+ *
+ * E8M0_1   = 0b01111111; => 2^(127-127) = 1
+ * E8M0_2   = 0b10000000; => 2^(128-127) = 2^1 = 2
+ * E8M0_3   = 0b10000010; => 2^(130-127) = 2^3 = 8
+ * E8M0_135 = 0b10000111; => 2^(135-127) = 2^8 = 256
+ * E8M0_142 = 0b10001110; => 2^(142-127) = 2^15 = 32768
+ * E8M0_MIN = 0b00000000; => 2^-127
+ * E8M0_MAX = 0b11111110; => 2^127
+ * E8M0_NAN = 0b11111111; => NaN
+ */
+struct e8m0_bexp_t
+{
+    using type = uint8_t;
+    type data;
+
+    constexpr static type bias     = 127;
+    constexpr static type nan_mask = 0xFF;
+
+    __host__ __device__ constexpr e8m0_bexp_t() : data{type{}} {}
+    __host__ __device__ constexpr e8m0_bexp_t(type init) : data{init} {}
+    __host__ __device__ constexpr e8m0_bexp_t(int init) : data{static_cast<type>(init & nan_mask)}
+    {
+    }
+    __host__ __device__ explicit constexpr e8m0_bexp_t(float scale)
+        : data{static_cast<type>((bit_cast<uint32_t>(scale) & (nan_mask << 23)) >> 23)}
+    {
+    }
+
+    __host__ __device__ explicit constexpr operator float() const
+    {
+        if(data == nan_mask || data == 0)
+        {
+            uint32_t bits = data << 1;
+            bits |= 1;
+            bits <<= 22;
+            return bit_cast<float>(bits);
+        }
+        else
+        {
+            uint32_t bits = data << 23;
+            return bit_cast<float>(bits);
+        }
+    }
+
+    __host__ __device__ constexpr bool operator==(const e8m0_bexp_t& other) const
+    {
+        // strict IEEE compliance for NaN
+        return data == other.data && data != nan_mask;
+    }
+
+    __host__ __device__ constexpr bool is_nan() const { return data == nan_mask; }
+};
+
+template <>
+struct numeric_traits<e8m0_bexp_t>
+{
+    static constexpr int PackedSize = 2;
+};
+
+namespace utils {
+
+template <typename T>
+__host__ __device__ inline constexpr int32_t get_exponent_value(T x);
+
+template <>
+__host__ __device__ inline constexpr int32_t get_exponent_value<e8m0_bexp_t>(e8m0_bexp_t x)
+{
+    return x.data;
+}
+
+} // namespace utils
+
+} // namespace ck_tile