diff --git a/hipamd/include/hip/amd_detail/amd_hip_fp8.h b/hipamd/include/hip/amd_detail/amd_hip_fp8.h
index e54c702417..47419cc61b 100644
--- a/hipamd/include/hip/amd_detail/amd_hip_fp8.h
+++ b/hipamd/include/hip/amd_detail/amd_hip_fp8.h
@@ -30,12 +30,28 @@
 #ifndef _HIP_INCLUDE_HIP_AMD_DETAIL_HIP_FP8_H_
 #define _HIP_INCLUDE_HIP_AMD_DETAIL_HIP_FP8_H_
 
-#if (defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)) && __HIP_DEVICE_COMPILE__
+#if (defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) || defined(__gfx1200__) ||  \
+     defined(__gfx1201__)) &&                                                                      \
+    __HIP_DEVICE_COMPILE__
 #define HIP_FP8_CVT_FAST_PATH 1
 #else
 #define HIP_FP8_CVT_FAST_PATH 0
 #endif
 
+#if (defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)) && __HIP_DEVICE_COMPILE__
+#define HIP_FP8_TYPE_OCP 0
+#define HIP_FP8_TYPE_FNUZ 1
+#elif (defined(__gfx1200__) || defined(__gfx1201__)) && __HIP_DEVICE_COMPILE__
+#define HIP_FP8_TYPE_OCP 1
+#define HIP_FP8_TYPE_FNUZ 0
+#elif __HIP_DEVICE_COMPILE__
+#define HIP_FP8_TYPE_FNUZ 0
+#define HIP_FP8_TYPE_OCP 0
+#else  // Host
+#define HIP_FP8_TYPE_FNUZ 1
+#define HIP_FP8_TYPE_OCP 1
+#endif
+
 #if !defined(__HIPCC_RTC__)
 #include <hip/amd_detail/amd_hip_common.h>
 #include <climits>
@@ -55,6 +71,8 @@
 #define __FP8_HOST_DEVICE_STATIC__ __FP8_HOST_DEVICE__ static inline
 #endif  // __HIPCC_RTC__
 
+#define __FP8_HOST__ __host__
+
 #if !defined(__HIPCC_RTC__)
 static_assert(CHAR_BIT == 8, "byte size should be of 8 bits");
 #endif
@@ -66,8 +84,10 @@ static_assert(sizeof(unsigned int) == 4);
  * \brief Describes FP8 interpretation
  */
 enum __hip_fp8_interpretation_t {
-  __HIP_E4M3_FNUZ = 0, /**< Standard FP8 */
-  __HIP_E5M2_FNUZ = 1, /**< BF8 */
+  __HIP_E4M3 = 0,      /**< OCP E4M3 */
+  __HIP_E5M2 = 1,      /**< OCP E5M2 */
+  __HIP_E4M3_FNUZ = 2, /**< Standard FP8*/
+  __HIP_E5M2_FNUZ = 3, /**< BF8 */
 };
 
 /**
@@ -99,13 +119,13 @@ typedef unsigned short int __hip_fp8x2_storage_t;
 typedef unsigned int __hip_fp8x4_storage_t;
 
 namespace internal {
+
 // The conversion function is from rocblas
 // https://github.com/ROCm/rocBLAS/blob/9b7f692abe3c54b88d1e77e045a7db7f1f188b69/library/include/internal/rocblas_hip_f8_impl.h#L39
 // This has been modified to add double types conversion as well
-template <typename T, bool negative_zero_nan>
+template <typename T, bool is_fnuz>
 __FP8_HOST_DEVICE_STATIC__ __hip_fp8_storage_t cast_to_f8(T _x, int wm, int we, bool clip = false,
-                                                          bool stoch = false,
-                                                          unsigned int rng = 0) {
+                                                   bool stoch = false, unsigned int rng = 0) {
   constexpr bool is_half = __hip_internal::is_same<T, _Float16>::value;
   constexpr bool is_float = __hip_internal::is_same<T, float>::value;
   constexpr bool is_double = __hip_internal::is_same<T, double>::value;
@@ -125,6 +145,7 @@ __FP8_HOST_DEVICE_STATIC__ __hip_fp8_storage_t cast_to_f8(T _x, int wm, int we,
   unsigned long long head, mantissa;
   int exponent, bias;
   unsigned int sign;
+  unsigned long long fInf, mask;
 
   if (sizeof(T) == 8) {
     head = x & 0xFFF0000000000000ull;
@@ -132,41 +153,80 @@ __FP8_HOST_DEVICE_STATIC__ __hip_fp8_storage_t cast_to_f8(T _x, int wm, int we,
     exponent = (head >> 52) & 0x7FF;
     sign = head >> 63;
     bias = 1023;
+    fInf = 0x7FF0000000000000ull;
+    mask = 0x7FFFFFFFFFFFFFFFull;
   } else if (sizeof(T) == 4) {
     head = x & 0xFF800000;
     mantissa = x & 0x7FFFFF;
     exponent = (head >> 23) & 0xFF;
     sign = head >> 31;
     bias = 127;
+    fInf = 0x7F800000;
+    mask = 0x7FFFFFFF;
   } else {
     head = x & 0xFC00;
     mantissa = x & 0x3FF;
     exponent = (head >> 10) & 0x1F;
     sign = head >> 15;
     bias = 15;
+    fInf = 0x7C00;
+    mask = 0x7FFF;
   }
-
-  unsigned int signed_inf = (sign << 7) + (((1 << we) - 1) << wm);
-
-  // Deal with inf and NaNs
-  if (negative_zero_nan) {
-    if (sizeof(T) == 8) {
-      if ((x & 0x7FF0000000000000ull) == 0x7FF0000000000000ull) return 0x80;
-    } else if (sizeof(T) == 4) {
-      if ((x & 0x7F800000) == 0x7F800000) return 0x80;
+  unsigned int signed_inf = 0;
+  unsigned int nan = 0;
+  if (is_fnuz) {
+    signed_inf = clip ? ((sign << 7) + 0x7f) : 0x80;
+    nan = 0x80;
+  } else {
+    if (we == 4) {  // e4m3
+      signed_inf = (sign << 7) + (clip ? 0x7e : 0x7f);
+    } else {  // e5m2
+      signed_inf = (sign << 7) + (clip ? 0x7b : 0x7c);
+    }
+    nan = (sign << 7) + 0x7f;
+  }
+  // Max values
+  unsigned long long ifmax = 0;
+  if (sizeof(T) == 8) {
+    if (we == 5) {  // 57344
+      ifmax = 0x40EC000000000000ull;
     } else {
-      if ((x & 0x7C00) == 0x7C00) return 0x80;
+      if (is_fnuz) {  // 240
+        ifmax = 0x406E000000000000ull;
+      } else {  // 448
+        ifmax = 0x407C000000000000ull;
+      }
+    }
+  } else if (sizeof(T) == 4) {
+    if (we == 5) {
+      ifmax = 0x47600000;
+    } else {
+      if (is_fnuz) {
+        ifmax = 0x43700000;
+      } else {
+        ifmax = 0x43E00000;
+      }
     }
   } else {
-    if (sizeof(T) == 8) {
-      if ((x & 0x7FF0000000000000ull) == 0x7FF0000000000000ull)
-        return signed_inf + (mantissa != 0 ? 1 : 0);
-    } else if (sizeof(T) == 4) {
-      if ((x & 0x7F800000) == 0x7F800000) return signed_inf + (mantissa != 0 ? 1 : 0);
+    if (we == 5) {
+      ifmax = 0x7B00;
     } else {
-      if ((x & 0x7C00) == 0x7C00) return signed_inf + (mantissa != 0 ? 1 : 0);
+      if (is_fnuz) {
+        ifmax = 0x5B80;
+      } else {
+        ifmax = 0x5F00;
+      }
     }
   }
+  // Deal with inf and NaNs
+  if ((x & fInf) == fInf) {
+    if (is_fnuz) return signed_inf;
+    return mantissa != 0 ? nan : signed_inf;
+  }
+
+  if ((x & mask) > ifmax) {
+    return signed_inf;
+  }
 
   if (x == 0) {
     return 0;
@@ -179,7 +239,7 @@ __FP8_HOST_DEVICE_STATIC__ __hip_fp8_storage_t cast_to_f8(T _x, int wm, int we,
   // exponent and mantissa again
 
   // For IEEE bias mode, the bias is 2^(k-1) -1 where k is the width of exponent bits
-  const int f8_bias = (1 << (we - 1)) - 1 + (negative_zero_nan ? 1 : 0);
+  const int f8_bias = (1 << (we - 1)) - 1 + (is_fnuz ? 1 : 0);
   const int f8_denormal_act_exponent = 1 - f8_bias;  // actual exponent of f8 denormal
   // act_exponent is the actual exponent of fp32/fp16 (after subtracting bias)
   // f8_exponent is the converted f8 exponent with bias encoding
@@ -252,7 +312,7 @@ after shift right by 4 bits, it would look like midpoint.
   mantissa >>= (mfmt - wm);
 
   // above range: quantize to maximum possible float of the same sign
-  const int max_exp = (1 << we) - (negative_zero_nan ? 1 : 2);
+  const int max_exp = (1 << we) - 1;
   if (f8_exponent > max_exp) {
     if (clip) {
       mantissa = (1 << wm) - 1;
@@ -262,16 +322,15 @@ after shift right by 4 bits, it would look like midpoint.
     }
   }
 
-  if (f8_exponent == 0 && mantissa == 0) return negative_zero_nan ? 0 : (sign << 7);
+  if (f8_exponent == 0 && mantissa == 0) return is_fnuz ? 0 : (sign << 7);
   mantissa &= (1 << wm) - 1;
   return (sign << 7) | (f8_exponent << wm) | mantissa;
 }
-
 // The conversion function is from rocblas
 // https://github.com/ROCm/rocBLAS/blob/9b7f692abe3c54b88d1e77e045a7db7f1f188b69/library/include/internal/rocblas_hip_f8_impl.h#L220
 // This has been modified to handle double types as well
-template <typename T, bool negative_zero_nan>
-__FP8_HOST_DEVICE_STATIC__ T cast_from_f8(__hip_fp8_storage_t x, int wm, int we) {
+template <typename T, bool is_fnuz>
+__FP8_HOST_DEVICE_STATIC__ T cast_from_f8(__hip_fp8_storage_t x, int wm, int we, bool clip = false) {
   constexpr bool is_half = __hip_internal::is_same<T, _Float16>::value;
   constexpr bool is_float = __hip_internal::is_same<T, float>::value;
   constexpr bool is_double = __hip_internal::is_same<T, double>::value;
@@ -280,34 +339,49 @@ __FP8_HOST_DEVICE_STATIC__ T cast_from_f8(__hip_fp8_storage_t x, int wm, int we)
   constexpr int weo = is_half ? 5 : (is_float ? 8 : 11);
   constexpr int wmo = is_half ? 10 : (is_float ? 23 : 52);
 
-  T fInf, fNegInf, fNaN, fNeg0;
+  T fInf, fNegInf, fNaN, fNeg0, fmax, fmin;
   if (is_half) {
     const unsigned short int ihInf = 0x7C00;
     const unsigned short int ihNegInf = 0xFC00;
     const unsigned short int ihNaN = 0x7C01;
     const unsigned short int ihNeg0 = 0x8000;
+    /* Max number in e5m2 57344*/
+    const unsigned short int ifmax = 0x7B00;
+    const unsigned short int ifmin = 0xFB00;
     fInf = reinterpret_cast<const _Float16&>(ihInf);
     fNegInf = reinterpret_cast<const _Float16&>(ihNegInf);
     fNaN = reinterpret_cast<const _Float16&>(ihNaN);
     fNeg0 = reinterpret_cast<const _Float16&>(ihNeg0);
+    fmax = reinterpret_cast<const _Float16&>(ifmax);
+    fmin = reinterpret_cast<const _Float16&>(ifmin);
   } else if (is_float) {
     const unsigned int ifInf = 0x7F800000;
     const unsigned int ifNegInf = 0xFF800000;
     const unsigned int ifNaN = 0x7F800001;
     const unsigned int ifNeg0 = 0x80000000;
+    /* Max number in e5m2 57344*/
+    const unsigned int ifmax = 0x47600000;
+    const unsigned int ifmin = 0xC7600000;
     fInf = reinterpret_cast<const float&>(ifInf);
     fNegInf = reinterpret_cast<const float&>(ifNegInf);
     fNaN = reinterpret_cast<const float&>(ifNaN);
     fNeg0 = reinterpret_cast<const float&>(ifNeg0);
+    fmax = reinterpret_cast<const float&>(ifmax);
+    fmin = reinterpret_cast<const float&>(ifmin);
   } else if (is_double) {
     const unsigned long long ifInf = 0x7FF0000000000000ull;
     const unsigned long long ifNegInf = 0xFFF0000000000000ull;
     const unsigned long long ifNaN = 0x7FF0000000000001ull;
     const unsigned long long ifNeg0 = 0x8000000000000000ull;
+    /* Max number in e5m2 57344*/
+    const unsigned long long ifmax = 0x40EC000000000000ull;
+    const unsigned long long ifmin = 0xC0EC000000000000ull;
     fInf = reinterpret_cast<const double&>(ifInf);
     fNegInf = reinterpret_cast<const double&>(ifNegInf);
     fNaN = reinterpret_cast<const double&>(ifNaN);
     fNeg0 = reinterpret_cast<const double&>(ifNeg0);
+    fmax = reinterpret_cast<const double&>(ifmax);
+    fmin = reinterpret_cast<const double&>(ifmin);
   }
 
   if (x == 0) {
@@ -317,11 +391,27 @@ __FP8_HOST_DEVICE_STATIC__ T cast_from_f8(__hip_fp8_storage_t x, int wm, int we)
   unsigned long long sign = x >> 7;
   unsigned long long mantissa = x & ((1 << wm) - 1);
   int exponent = (x & 0x7F) >> wm;
-  if (negative_zero_nan) {
-    if (x == 0x80) return fNaN;
+  if (is_fnuz) {
+    if (x == 0x80) {
+      return fNaN;
+    }
   } else {
-    if (x == 0x80) return fNeg0;
-    if (exponent == ((1 << we) - 1)) return (mantissa == 0) ? (sign ? fNegInf : fInf) : fNaN;
+    if (x == 0x80) {
+      return fNeg0;
+    }
+    if (we == 4) {  // e4m3
+      if ((x & 0x7F) == 0x7F) {
+        return fNaN;
+      }
+    } else if ((x & 0x7C) == 0x7C) {  // e5m2
+      if ((x & 0x3) == 0) {
+        if (clip) {
+          return sign ? fmin : fmax;
+        }
+        return sign ? fNegInf : fInf;
+      }
+      return fNaN;
+    }
   }
 
   typename __hip_internal::conditional<
@@ -329,12 +419,12 @@ __FP8_HOST_DEVICE_STATIC__ T cast_from_f8(__hip_fp8_storage_t x, int wm, int we)
       typename __hip_internal::conditional<sizeof(T) == 4, unsigned int,
                                            unsigned long long>::type>::type retval;
 
-  if (we == 5 && is_half && !negative_zero_nan) {
+  if (we == 5 && is_half && !is_fnuz) {
     retval = x << 8;
     return reinterpret_cast<const T&>(retval);
   }
 
-  const int exp_low_cutoff = (1 << (weo - 1)) - (1 << (we - 1)) + 1 - (negative_zero_nan ? 1 : 0);
+  const int exp_low_cutoff = (1 << (weo - 1)) - (1 << (we - 1)) + 1 - (is_fnuz ? 1 : 0);
 
   // subnormal input
   if (exponent == 0) {
@@ -389,6 +479,10 @@ static __device__ __hip_fp8_storage_t cast_to_f8_from_f32(float v, bool saturate
       if ((val.i32val & 0x7F800000) != 0x7F800000) {  /// propagate NAN/INF, no clipping
         val.fval = __builtin_amdgcn_fmed3f(val.fval, 240.0, -240.0);
       }
+    } else if (interpret == __HIP_E4M3) {             // OCP type
+      if ((val.i32val & 0x7F800000) != 0x7F800000) {  /// propagate NAN/INF, no clipping
+        val.fval = __builtin_amdgcn_fmed3f(val.fval, 448.0, -448.0);
+      }
     } else {
       if ((val.i32val & 0x7F800000) != 0x7F800000) {  /// propagate NAN/INF, no clipping
         val.fval = __builtin_amdgcn_fmed3f(val.fval, 57344.0, -57344.0);
@@ -397,13 +491,13 @@ static __device__ __hip_fp8_storage_t cast_to_f8_from_f32(float v, bool saturate
   }
 
   if (stochastic_rounding) {
-    ival = interpret == __HIP_E4M3_FNUZ
+    ival = (interpret == __HIP_E4M3_FNUZ) || (interpret == __HIP_E4M3)
         ? __builtin_amdgcn_cvt_sr_fp8_f32(val.fval, rng, ival, 0)
         : __builtin_amdgcn_cvt_sr_bf8_f32(val.fval, rng, ival, 0);  // 0 pos
     val.i32val = ival;
     i8data = val.i8val[0];  // little endian
   } else {                  // RNE CVT
-    ival = interpret == __HIP_E4M3_FNUZ
+    ival = (interpret == __HIP_E4M3_FNUZ) || (interpret == __HIP_E4M3)
         ? __builtin_amdgcn_cvt_pk_fp8_f32(val.fval, val.fval, ival, false)
         : __builtin_amdgcn_cvt_pk_bf8_f32(val.fval, val.fval, ival, false);  // false -> WORD0
     val.i32val = ival;
@@ -425,15 +519,31 @@ cast_to_f8x2_from_f32x2(float2 v, bool saturate, __hip_fp8_interpretation_t inte
   f2val.fval = v;
 
   if (saturate) {  /// propagate NAN/INF, no clipping
-    if ((f2val.i32val[0] & 0x7F800000) != 0x7F800000) {
-      f2val.fval.x = __builtin_amdgcn_fmed3f(f2val.fval.x, 240.0, -240.0);
-    }
-    if ((f2val.i32val[1] & 0x7F800000) != 0x7F800000) {
-      f2val.fval.y = __builtin_amdgcn_fmed3f(f2val.fval.x, 240.0, -240.0);
+    if (interpret == __HIP_E4M3_FNUZ) {
+      if ((f2val.i32val[0] & 0x7F800000) != 0x7F800000) {
+        f2val.fval.x = __builtin_amdgcn_fmed3f(f2val.fval.x, 240.0, -240.0);
+      }
+      if ((f2val.i32val[1] & 0x7F800000) != 0x7F800000) {
+        f2val.fval.y = __builtin_amdgcn_fmed3f(f2val.fval.x, 240.0, -240.0);
+      }
+    } else if (interpret == __HIP_E4M3) {
+      if ((f2val.i32val[0] & 0x7F800000) != 0x7F800000) {
+        f2val.fval.x = __builtin_amdgcn_fmed3f(f2val.fval.x, 448.0, -448.0);
+      }
+      if ((f2val.i32val[1] & 0x7F800000) != 0x7F800000) {
+        f2val.fval.y = __builtin_amdgcn_fmed3f(f2val.fval.x, 448.0, -448.0);
+      }
+    } else {
+      if ((f2val.i32val[0] & 0x7F800000) != 0x7F800000) {
+        f2val.fval.x = __builtin_amdgcn_fmed3f(f2val.fval.x, 57344.0, -57344.0);
+      }
+      if ((f2val.i32val[1] & 0x7F800000) != 0x7F800000) {
+        f2val.fval.y = __builtin_amdgcn_fmed3f(f2val.fval.x, 57344.0, -57344.0);
+      }
     }
   }
 
-  f2val.i32val[0] = interpret == __HIP_E4M3_FNUZ
+  f2val.i32val[0] = (interpret == __HIP_E4M3_FNUZ) || (interpret == __HIP_E4M3)
       ? __builtin_amdgcn_cvt_pk_fp8_f32(v.x, v.y, 0, false)
       : __builtin_amdgcn_cvt_pk_bf8_f32(v.x, v.y, 0, false);
 
@@ -448,8 +558,9 @@ static __device__ float cast_to_f32_from_f8(__hip_fp8_storage_t v,
   } val;
   val.i8val[0] = v;
 
-  float fval = interpret == __HIP_E4M3_FNUZ ? __builtin_amdgcn_cvt_f32_fp8(val.i32val, 0)
-                                            : __builtin_amdgcn_cvt_f32_bf8(val.i32val, 0);
+  float fval = (interpret == __HIP_E4M3_FNUZ) || (interpret == __HIP_E4M3)
+      ? __builtin_amdgcn_cvt_f32_fp8(val.i32val, 0)
+      : __builtin_amdgcn_cvt_f32_bf8(val.i32val, 0);
   return fval;
 }
 
@@ -461,8 +572,9 @@ static __device__ float2 cast_to_f32x2_from_f8x2(__hip_fp8x2_storage_t v,
   } val;
   val.i16val[0] = v;
 
-  auto f2 = interpret == __HIP_E4M3_FNUZ ? __builtin_amdgcn_cvt_pk_f32_fp8(val.i32val, false)
-                                         : __builtin_amdgcn_cvt_pk_f32_bf8(val.i32val, false);
+  auto f2 = (interpret == __HIP_E4M3_FNUZ) || (interpret == __HIP_E4M3)
+      ? __builtin_amdgcn_cvt_pk_f32_fp8(val.i32val, false)
+      : __builtin_amdgcn_cvt_pk_f32_bf8(val.i32val, false);
   return float2{f2[0], f2[1]};
 }
 #endif  // HIP_FP8_CVT_FAST_PATH
@@ -473,6 +585,19 @@ NaN are represented by 1-0000-000 or 1-00000-00 */
 __FP8_HOST_DEVICE_STATIC__ bool hip_fp8_fnuz_is_nan(__hip_fp8_storage_t a) {
   return static_cast<unsigned char>(a) == 0x80;
 }
+
+__FP8_HOST_DEVICE_STATIC__ bool hip_fp8_ocp_is_nan(__hip_fp8_storage_t a,
+                                                   const __hip_fp8_interpretation_t type) {
+  return (type == __HIP_E4M3) ? ((a & 0x7f) == 0x7f)
+      : (type == __HIP_E5M2)  ? ((a & 0x7f) > 0x7c)
+                              : false;
+}
+
+__FP8_HOST_DEVICE_STATIC__ bool hip_fp8_ocp_is_inf(__hip_fp8_storage_t a,
+                                                   const __hip_fp8_interpretation_t type) {
+  return (type == __HIP_E5M2) ? (a & 0x7f) == 0x7c : false;
+}
+
 }  // namespace internal
 
 /**
@@ -488,12 +613,20 @@ __FP8_HOST_DEVICE_STATIC__ __hip_fp8_storage_t __hip_cvt_float_to_fp8(
 #if HIP_FP8_CVT_FAST_PATH
   return internal::cast_to_f8_from_f32<false>(f, sat == __HIP_SATFINITE, type);
 #else   // HIP_FP8_CVT_FAST_PATH
-  int we = type == __HIP_E4M3_FNUZ ? 4 : 5;
-  int wm = type == __HIP_E4M3_FNUZ ? 3 : 2;
-  return internal::cast_to_f8<float, true>(f, wm, we, sat == __HIP_SATFINITE);
+  if (type == __HIP_E4M3_FNUZ || type == __HIP_E5M2_FNUZ) {
+    int we = type == __HIP_E4M3_FNUZ ? 4 : 5;
+    int wm = type == __HIP_E4M3_FNUZ ? 3 : 2;
+    return internal::cast_to_f8<float, true>(f, wm, we, sat == __HIP_SATFINITE);
+  }
+  if (type == __HIP_E4M3 || type == __HIP_E5M2) {
+    int we = type == __HIP_E4M3 ? 4 : 5;
+    int wm = type == __HIP_E4M3 ? 3 : 2;
+    return internal::cast_to_f8<float, false>(f, wm, we, sat == __HIP_SATFINITE);
+  }
 #endif  // HIP_FP8_CVT_FAST_PATH
 }
 
+
 /**
  * \brief convert float2 to @p __hip_fp8x2_storage_t
  *
@@ -523,9 +656,16 @@ __FP8_HOST_DEVICE_STATIC__ __hip_fp8x2_storage_t __hip_cvt_float2_to_fp8x2(
  */
 __FP8_HOST_DEVICE_STATIC__ __hip_fp8_storage_t __hip_cvt_double_to_fp8(
     const double d, const __hip_saturation_t sat, const __hip_fp8_interpretation_t type) {
-  int we = type == __HIP_E4M3_FNUZ ? 4 : 5;
-  int wm = type == __HIP_E4M3_FNUZ ? 3 : 2;
-  return internal::cast_to_f8<double, true>(d, wm, we, sat == __HIP_SATFINITE);
+  if (type == __HIP_E4M3_FNUZ || type == __HIP_E5M2_FNUZ) {
+    int we = type == __HIP_E4M3_FNUZ ? 4 : 5;
+    int wm = type == __HIP_E4M3_FNUZ ? 3 : 2;
+    return internal::cast_to_f8<double, true>(d, wm, we, sat == __HIP_SATFINITE);
+  }
+  if (type == __HIP_E4M3 || type == __HIP_E5M2) {
+    int we = type == __HIP_E4M3 ? 4 : 5;
+    int wm = type == __HIP_E4M3 ? 3 : 2;
+    return internal::cast_to_f8<double, false>(d, wm, we, sat == __HIP_SATFINITE);
+  }
 }
 
 /**
@@ -582,9 +722,16 @@ __hip_cvt_bfloat16raw2_to_fp8x2(const __hip_bfloat162_raw hr, const __hip_satura
  */
 __FP8_HOST_DEVICE_STATIC__ __half_raw
 __hip_cvt_fp8_to_halfraw(const __hip_fp8_storage_t x, const __hip_fp8_interpretation_t type) {
-  unsigned int we = type == __HIP_E4M3_FNUZ ? 4 : 5;
-  unsigned int wm = type == __HIP_E4M3_FNUZ ? 3 : 2;
-  return __half_raw{internal::cast_from_f8<_Float16, true>(x, wm, we)};
+  if (type == __HIP_E4M3_FNUZ || type == __HIP_E5M2_FNUZ) {
+    unsigned int we = type == __HIP_E4M3_FNUZ ? 4 : 5;
+    unsigned int wm = type == __HIP_E4M3_FNUZ ? 3 : 2;
+    return __half_raw{internal::cast_from_f8<_Float16, true>(x, wm, we)};
+  }
+  if (type == __HIP_E4M3 || type == __HIP_E5M2) {
+    unsigned int we = type == __HIP_E4M3 ? 4 : 5;
+    unsigned int wm = type == __HIP_E4M3 ? 3 : 2;
+    return __half_raw{internal::cast_from_f8<_Float16, false>(x, wm, we)};
+  }
 }
 
 /**
@@ -644,75 +791,135 @@ struct __hip_fp8_e4m3_fnuz {
   // Add cast from unsigned long long, long long to fp8
 
   /*! create fp8 e4m3 from long */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8_e4m3_fnuz(const long int val)
+#else
+  __FP8_HOST__ __hip_fp8_e4m3_fnuz(const long int val)
+#endif
       : __x(__hip_cvt_float_to_fp8(static_cast<float>(val), __default_saturation,
                                    __default_interpret)) {}
 
   /*! create fp8 e4m3 from int */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8_e4m3_fnuz(const int val)
+#else
+  __FP8_HOST__ __hip_fp8_e4m3_fnuz(const int val)
+#endif
       : __x(__hip_cvt_float_to_fp8(static_cast<float>(val), __default_saturation,
                                    __default_interpret)) {}
 
   /*! create fp8 e4m3 from short int */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8_e4m3_fnuz(const short int val)
+#else
+  __FP8_HOST__ __hip_fp8_e4m3_fnuz(const short int val)
+#endif
       : __x(__hip_cvt_float_to_fp8(static_cast<float>(val), __default_saturation,
                                    __default_interpret)) {}
 
   /*! create fp8 e4m3 from unsigned long */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8_e4m3_fnuz(const unsigned long int val)
+#else
+  __FP8_HOST__ __hip_fp8_e4m3_fnuz(const unsigned long int val)
+#endif
       : __x(__hip_cvt_float_to_fp8(static_cast<float>(val), __default_saturation,
                                    __default_interpret)) {}
 
   /*! create fp8 e4m3 from unsigned int */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8_e4m3_fnuz(const unsigned int val)
+#else
+  __FP8_HOST__ __hip_fp8_e4m3_fnuz(const unsigned int val)
+#endif
       : __x(__hip_cvt_float_to_fp8(static_cast<float>(val), __default_saturation,
                                    __default_interpret)) {}
 
   /*! create fp8 e4m3 from unsigned short */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8_e4m3_fnuz(const unsigned short int val)
+#else
+  __FP8_HOST__ __hip_fp8_e4m3_fnuz(const unsigned short int val)
+#endif
       : __x(__hip_cvt_float_to_fp8(static_cast<float>(val), __default_saturation,
                                    __default_interpret)) {}
 
   /*! create fp8 e4m3 from double */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8_e4m3_fnuz(const double f)
+#else
+  __FP8_HOST__ __hip_fp8_e4m3_fnuz(const double f)
+#endif
       : __x(__hip_cvt_double_to_fp8(f, __default_saturation, __default_interpret)) {}
 
   /*! create fp8 e4m3 from float */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8_e4m3_fnuz(const float f)
+#else
+  __FP8_HOST__ __hip_fp8_e4m3_fnuz(const float f)
+#endif
       : __x(__hip_cvt_float_to_fp8(f, __default_saturation, __default_interpret)) {}
 
   /*! create fp8 e4m3 from __hip_bfloat16 */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8_e4m3_fnuz(const __hip_bfloat16 f)
+#else
+  __FP8_HOST__ __hip_fp8_e4m3_fnuz(const __hip_bfloat16 f)
+#endif
       : __x(__hip_cvt_float_to_fp8(static_cast<float>(f), __default_saturation,
                                    __default_interpret)) {}
 
   /*! create fp8 e4m3 from __half */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8_e4m3_fnuz(const __half f)
+#else
+  __FP8_HOST__ __hip_fp8_e4m3_fnuz(const __half f)
+#endif
       : __x(__hip_cvt_halfraw_to_fp8(static_cast<__half_raw>(f), __default_saturation,
                                      __default_interpret)) {}
 
   /*! default construct fp8 e4m3 */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8_e4m3_fnuz() = default;
+#else
+  __FP8_HOST__ __hip_fp8_e4m3_fnuz() = default;
+#endif
 
   /*! convert fp8 e4m3 to __half */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator __half() const {
+#else
+  __FP8_HOST__ operator __half() const {
+#endif
     return __half(__hip_cvt_fp8_to_halfraw(__x, __default_interpret));
   }
 
   /*! convert fp8 e4m3 to __hip_bfloat16 */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator __hip_bfloat16() const {
+#else
+  __FP8_HOST__ operator __hip_bfloat16() const {
+#endif
     float f = *this;
     return __hip_bfloat16(f);
   }
 
   /*! convert fp8 e4m3 to bool, return false if value is 0, true otherwise */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator bool() const {
+#else
+  __FP8_HOST__ operator bool() const {
+#endif
     // it can be 0x00 (+0.0) since 0x80 will be nan
     return !(static_cast<unsigned short>(__x) == 0);
   }
 
   /*! convert fp8 e4m3 to char, clamp number to CHAR_MIN/CHAR_MAX if its out of range */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator char() const {
+#else
+  __FP8_HOST__ operator char() const {
+#endif
     if (internal::hip_fp8_fnuz_is_nan(__x)) {
       return 0;
     }
@@ -728,12 +935,20 @@ struct __hip_fp8_e4m3_fnuz {
   }
 
   /*! convert fp8 e4m3 to double */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator double() const {
+#else
+  __FP8_HOST__ operator double() const {
+#endif
     return internal::cast_from_f8<double, true>(__x, __wm, __we);
   }
 
   /*! convert fp8 e4m3 to float */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator float() const {
+#else
+  __FP8_HOST__ operator float() const {
+#endif
 #if HIP_FP8_CVT_FAST_PATH
     return internal::cast_to_f32_from_f8(__x, __default_interpret);
 #else
@@ -742,7 +957,11 @@ struct __hip_fp8_e4m3_fnuz {
   }
 
   /*! convert fp8 e4m3 to int, return 0 if value is NaN */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator int() const {
+#else
+  __FP8_HOST__ operator int() const {
+#endif
     if (internal::hip_fp8_fnuz_is_nan(__x)) {
       return 0;
     }
@@ -752,7 +971,11 @@ struct __hip_fp8_e4m3_fnuz {
   }
 
   /*! convert fp8 e4m3 to long, return 0 if value is NaN */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator long int() const {
+#else
+  __FP8_HOST__ operator long int() const {
+#endif
     if (internal::hip_fp8_fnuz_is_nan(__x)) {
       return 0;
     }
@@ -762,7 +985,11 @@ struct __hip_fp8_e4m3_fnuz {
   }
 
   /*! convert fp8 e4m3 to long long, return 0 if value is NaN */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator long long int() const {
+#else
+  __FP8_HOST__ operator long long int() const {
+#endif
     if (internal::hip_fp8_fnuz_is_nan(__x)) {
       return 0;
     }
@@ -772,7 +999,11 @@ struct __hip_fp8_e4m3_fnuz {
   }
 
   /*! convert fp8 e4m3 to short int, clamp out of bound values, return 0 if value is NaN */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator short int() const {
+#else
+  __FP8_HOST__ operator short int() const {
+#endif
     if (internal::hip_fp8_fnuz_is_nan(__x)) {
       return 0;
     }
@@ -788,7 +1019,11 @@ struct __hip_fp8_e4m3_fnuz {
   }
 
   /*! convert fp8 e4m3 to signed char, clamp out of bound values, return 0 if value is NaN */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator signed char() const {
+#else
+  __FP8_HOST__ operator signed char() const {
+#endif
     if (internal::hip_fp8_fnuz_is_nan(__x)) {
       return 0;
     }
@@ -804,7 +1039,11 @@ struct __hip_fp8_e4m3_fnuz {
   }
 
   /*! convert fp8 e4m3 to unsigned char, clamp out of bound values, return 0 if value is NaN */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator unsigned char() const {
+#else
+  __FP8_HOST__ operator unsigned char() const {
+#endif
     if (internal::hip_fp8_fnuz_is_nan(__x)) {
       return 0;
     }
@@ -820,7 +1059,11 @@ struct __hip_fp8_e4m3_fnuz {
   }
 
   /*! convert fp8 e4m3 to unsigned int, return 0 if value is NaN */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator unsigned int() const {
+#else
+  __FP8_HOST__ operator unsigned int() const {
+#endif
     if (internal::hip_fp8_fnuz_is_nan(__x)) {
       return 0;
     }
@@ -834,7 +1077,11 @@ struct __hip_fp8_e4m3_fnuz {
   }
 
   /*! convert fp8 e4m3 to unsigned long, return 0 if value is NaN */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator unsigned long int() const {
+#else
+  __FP8_HOST__ operator unsigned long int() const {
+#endif
     if (internal::hip_fp8_fnuz_is_nan(__x)) {
       return 0;
     }
@@ -848,7 +1095,11 @@ struct __hip_fp8_e4m3_fnuz {
   }
 
   /*! convert fp8 e4m3 to long long int, return 0 if value is NaN */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator unsigned long long int() const {
+#else
+  __FP8_HOST__ operator unsigned long long int() const {
+#endif
     if (internal::hip_fp8_fnuz_is_nan(__x)) {
       return 0;
     }
@@ -862,7 +1113,11 @@ struct __hip_fp8_e4m3_fnuz {
   }
 
   /*! convert fp8 e4m3 to unsigned short, return 0 if value is NaN */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator unsigned short int() const {
+#else
+  __FP8_HOST__ operator unsigned short int() const {
+#endif
     if (internal::hip_fp8_fnuz_is_nan(__x)) {
       return 0;
     }
@@ -888,31 +1143,59 @@ struct __hip_fp8x2_e4m3_fnuz {
   static constexpr unsigned int __wm = 3;
 
   /*! create fp8x2 e4m3 type from double2 */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8x2_e4m3_fnuz(const double2 val)
+#else
+  __FP8_HOST__ __hip_fp8x2_e4m3_fnuz(const double2 val)
+#endif
       : __x(__hip_cvt_double2_to_fp8x2(val, __default_saturation, __default_interpret)) {}
 
   /*! create fp8x2 e4m3 type from float2 */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8x2_e4m3_fnuz(const float2 val)
+#else
+  __FP8_HOST__ __hip_fp8x2_e4m3_fnuz(const float2 val)
+#endif
       : __x(__hip_cvt_float2_to_fp8x2(val, __default_saturation, __default_interpret)) {}
 
   /*! create fp8x2 e4m3 type from __hip_bfloat162 */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8x2_e4m3_fnuz(const __hip_bfloat162 val)
+#else
+  __FP8_HOST__ __hip_fp8x2_e4m3_fnuz(const __hip_bfloat162 val)
+#endif
       : __x(__hip_cvt_bfloat16raw2_to_fp8x2(val, __default_saturation, __default_interpret)) {}
 
   /*! create fp8x2 e4m3 type from __half2 */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8x2_e4m3_fnuz(const __half2 val)
+#else
+  __FP8_HOST__ __hip_fp8x2_e4m3_fnuz(const __half2 val)
+#endif
       : __x(__hip_cvt_halfraw2_to_fp8x2(val, __default_saturation, __default_interpret)) {}
 
   /*! Default construct of fp8x2 e4m3 */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8x2_e4m3_fnuz() = default;
+#else
+  __FP8_HOST__ __hip_fp8x2_e4m3_fnuz() = default;
+#endif
 
   /*! convert fp8x2 e4m3 to __half2 */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator __half2() const {
+#else
+  __FP8_HOST__ operator __half2() const {
+#endif
     return __half2(__hip_cvt_fp8x2_to_halfraw2(__x, __default_interpret));
   }
 
   /*! convert fp8x2 e4m3 to float2 */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator float2() const {
+#else
+  __FP8_HOST__ operator float2() const {
+#endif
 #if HIP_FP8_CVT_FAST_PATH
     return internal::cast_to_f32x2_from_f8x2(__x, __default_interpret);
 #else
@@ -936,7 +1219,11 @@ struct __hip_fp8x4_e4m3_fnuz {
   static constexpr unsigned int __wm = 3;
 
   /*! create fp8x4 e4m3 type from double4 */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8x4_e4m3_fnuz(const double4 val)
+#else
+  __FP8_HOST__ __hip_fp8x4_e4m3_fnuz(const double4 val)
+#endif
       : __x{reinterpret_cast<__hip_fp8x4_storage_t>(
             static_cast<unsigned int>(reinterpret_cast<unsigned char>(__hip_cvt_double_to_fp8(
                                           val.x, __default_saturation, __default_interpret)) |
@@ -951,7 +1238,11 @@ struct __hip_fp8x4_e4m3_fnuz {
                                           << 24))} {}
 
   /*! create fp8x4 e4m3 type from float4 */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8x4_e4m3_fnuz(const float4 val)
+#else
+  __FP8_HOST__ __hip_fp8x4_e4m3_fnuz(const float4 val)
+#endif
       : __x{reinterpret_cast<__hip_fp8x4_storage_t>(
             static_cast<unsigned int>(reinterpret_cast<unsigned char>(__hip_cvt_float_to_fp8(
                                           val.x, __default_saturation, __default_interpret)) |
@@ -966,7 +1257,11 @@ struct __hip_fp8x4_e4m3_fnuz {
                                           << 24))} {}
 
   /*! create fp8x4 e4m3 type from two __hip_bfloat162 */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8x4_e4m3_fnuz(const __hip_bfloat162 low, const __hip_bfloat162 high)
+#else
+  __FP8_HOST__ __hip_fp8x4_e4m3_fnuz(const __hip_bfloat162 low, const __hip_bfloat162 high)
+#endif
       : __x(reinterpret_cast<__hip_fp8x4_storage_t>(static_cast<unsigned int>(
             reinterpret_cast<unsigned short>(
                 __hip_cvt_bfloat16raw2_to_fp8x2(high, __default_saturation, __default_interpret)) |
@@ -975,7 +1270,11 @@ struct __hip_fp8x4_e4m3_fnuz {
                 << 16))) {}
 
   /*! create fp8x4 e4m3 type from two __half2 */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8x4_e4m3_fnuz(const __half2 low, const __half2 high)
+#else
+  __FP8_HOST__ __hip_fp8x4_e4m3_fnuz(const __half2 low, const __half2 high)
+#endif
       : __x(reinterpret_cast<__hip_fp8x4_storage_t>(
             static_cast<unsigned int>(reinterpret_cast<unsigned short>(__hip_cvt_halfraw2_to_fp8x2(
                                           high, __default_saturation, __default_interpret)) |
@@ -984,10 +1283,18 @@ struct __hip_fp8x4_e4m3_fnuz {
                                           << 16))) {}
 
   /*! Default construct fp8x4 e4m3 */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8x4_e4m3_fnuz() = default;
+#else
+  __FP8_HOST__ __hip_fp8x4_e4m3_fnuz() = default;
+#endif
 
   /*! convert fp8x4 e4m3 to float4 */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator float4() const {
+#else
+  __FP8_HOST__ operator float4() const {
+#endif
     auto x = __x;                                                    // bypass const
     auto fp8x2_low = *reinterpret_cast<__hip_fp8x2_storage_t*>(&x);  // Little E
     auto fp8x2_high = *(reinterpret_cast<__hip_fp8x2_storage_t*>(&x) + 1);
@@ -1024,58 +1331,106 @@ struct __hip_fp8_e5m2_fnuz {
   // Add cast from unsigned long long, long long to fp8
 
   /*! create fp8 e5m2 type from long */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8_e5m2_fnuz(const long int val)
+#else
+  __FP8_HOST__ __hip_fp8_e5m2_fnuz(const long int val)
+#endif
       : __x(__hip_cvt_float_to_fp8(static_cast<float>(val), __default_saturation,
                                    __default_interpret)) {}
 
   /*! create fp8 e5m2 type from int */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8_e5m2_fnuz(const int val)
+#else
+  __FP8_HOST__ __hip_fp8_e5m2_fnuz(const int val)
+#endif
       : __x(__hip_cvt_float_to_fp8(static_cast<float>(val), __default_saturation,
                                    __default_interpret)) {}
 
   /*! create fp8 e5m2 type from short int */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8_e5m2_fnuz(const short int val)
+#else
+  __FP8_HOST__ __hip_fp8_e5m2_fnuz(const short int val)
+#endif
       : __x(__hip_cvt_float_to_fp8(static_cast<float>(val), __default_saturation,
                                    __default_interpret)) {}
 
   /*! create fp8 e5m2 type from unsigned long */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8_e5m2_fnuz(const unsigned long int val)
+#else
+  __FP8_HOST__ __hip_fp8_e5m2_fnuz(const unsigned long int val)
+#endif
       : __x(__hip_cvt_float_to_fp8(static_cast<float>(val), __default_saturation,
                                    __default_interpret)) {}
 
   /*! create fp8 e5m2 type from unsigned int */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8_e5m2_fnuz(const unsigned int val)
+#else
+  __FP8_HOST__ __hip_fp8_e5m2_fnuz(const unsigned int val)
+#endif
       : __x(__hip_cvt_float_to_fp8(static_cast<float>(val), __default_saturation,
                                    __default_interpret)) {}
 
   /*! create fp8 e5m2 type from unsigned short */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8_e5m2_fnuz(const unsigned short int val)
+#else
+  __FP8_HOST__ __hip_fp8_e5m2_fnuz(const unsigned short int val)
+#endif
       : __x(__hip_cvt_float_to_fp8(static_cast<float>(val), __default_saturation,
                                    __default_interpret)) {}
 
   /*! create fp8 e5m2 type from double */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8_e5m2_fnuz(const double f)
+#else
+  __FP8_HOST__ __hip_fp8_e5m2_fnuz(const double f)
+#endif
       : __x(__hip_cvt_double_to_fp8(f, __default_saturation, __default_interpret)) {}
 
   /*! create fp8 e5m2 type from float */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8_e5m2_fnuz(const float f)
+#else
+  __FP8_HOST__ __hip_fp8_e5m2_fnuz(const float f)
+#endif
       : __x(__hip_cvt_float_to_fp8(f, __default_saturation, __default_interpret)) {}
 
   /*! create fp8 e5m2 type from __hip_bfloat16 */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8_e5m2_fnuz(const __hip_bfloat16 f)
+#else
+  __FP8_HOST__ __hip_fp8_e5m2_fnuz(const __hip_bfloat16 f)
+#endif
       : __x(__hip_cvt_float_to_fp8(static_cast<float>(f), __default_saturation,
                                    __default_interpret)) {}
 
   /*! create fp8 e5m2 type from __hip_bfloat16 */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8_e5m2_fnuz(const __half f)
+#else
+  __FP8_HOST__ __hip_fp8_e5m2_fnuz(const __half f)
+#endif
       : __x(__hip_cvt_halfraw_to_fp8(static_cast<__half_raw>(f), __default_saturation,
                                      __default_interpret)) {}
 
   /*! default construct fp8 e5m2 */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8_e5m2_fnuz() = default;
+#else
+  __FP8_HOST__ __hip_fp8_e5m2_fnuz() = default;
+#endif
 
   /*! convert fp8 e5m2 to float */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator float() const {
+#else
+  __FP8_HOST__ operator float() const {
+#endif
 #if HIP_FP8_CVT_FAST_PATH
     return internal::cast_to_f32_from_f8(__x, __default_interpret);
 #else
@@ -1084,24 +1439,40 @@ struct __hip_fp8_e5m2_fnuz {
   }
 
   /*! convert fp8 e5m2 to __half */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator __half() const {
+#else
+  __FP8_HOST__ operator __half() const {
+#endif
     return __half(__hip_cvt_fp8_to_halfraw(__x, __default_interpret));
   }
 
   /*! convert fp8 e5m2 to __hip_bfloat16 */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator __hip_bfloat16() const {
+#else
+  __FP8_HOST__ operator __hip_bfloat16() const {
+#endif
     float f = *this;
     return __hip_bfloat16(f);
   }
 
   /*! convert fp8 e4m3 to bool, return false if value is 0, true otherwise */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator bool() const {
+#else
+  __FP8_HOST__ operator bool() const {
+#endif
     // it can be 0x00 (+0.0) since 0x80 will be nan
     return !(static_cast<unsigned short>(__x) == 0);
   }
 
   /*! convert fp8 e5m2 to char, clamp out of bound values, return 0 if value is NaN */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator char() const {
+#else
+  __FP8_HOST__ operator char() const {
+#endif
     if (internal::hip_fp8_fnuz_is_nan(__x)) {
       return 0;
     }
@@ -1117,12 +1488,20 @@ struct __hip_fp8_e5m2_fnuz {
   }
 
   /*! convert fp8 e5m2 to double */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator double() const {
+#else
+  __FP8_HOST__ operator double() const {
+#endif
     return internal::cast_from_f8<double, true>(__x, __wm, __we);
   }
 
   /*! convert fp8 e5m2 to int, return 0 if value is NaN */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator int() const {
+#else
+  __FP8_HOST__ operator int() const {
+#endif
     if (internal::hip_fp8_fnuz_is_nan(__x)) {
       return 0;
     }
@@ -1132,7 +1511,11 @@ struct __hip_fp8_e5m2_fnuz {
   }
 
   /*! convert fp8 e5m2 to long, return 0 if value is NaN */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator long int() const {
+#else
+  __FP8_HOST__ operator long int() const {
+#endif
     if (internal::hip_fp8_fnuz_is_nan(__x)) {
       return 0;
     }
@@ -1142,7 +1525,11 @@ struct __hip_fp8_e5m2_fnuz {
   }
 
   /*! convert fp8 e5m2 to long long, return 0 if value is NaN */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator long long int() const {
+#else
+  __FP8_HOST__ operator long long int() const {
+#endif
     if (internal::hip_fp8_fnuz_is_nan(__x)) {
       return 0;
     }
@@ -1152,7 +1539,11 @@ struct __hip_fp8_e5m2_fnuz {
   }
 
   /*! convert fp8 e5m2 to short, clamp out of bound values, return 0 if value is NaN */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator short int() const {
+#else
+  __FP8_HOST__ operator short int() const {
+#endif
     if (internal::hip_fp8_fnuz_is_nan(__x)) {
       return 0;
     }
@@ -1168,7 +1559,11 @@ struct __hip_fp8_e5m2_fnuz {
   }
 
   /*! convert fp8 e5m2 to signed char, clamp out of bound values, return 0 if value is NaN */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator signed char() const {
+#else
+  __FP8_HOST__ operator signed char() const {
+#endif
     if (internal::hip_fp8_fnuz_is_nan(__x)) {
       return 0;
     }
@@ -1184,7 +1579,11 @@ struct __hip_fp8_e5m2_fnuz {
   }
 
   /*! convert fp8 e5m2 to unsigned char, clamp out of bound values, return 0 if value is NaN */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator unsigned char() const {
+#else
+  __FP8_HOST__ operator unsigned char() const {
+#endif
     if (internal::hip_fp8_fnuz_is_nan(__x)) {
       return 0;
     }
@@ -1200,7 +1599,11 @@ struct __hip_fp8_e5m2_fnuz {
   }
 
   /*! convert fp8 e5m2 to unsigned int, return 0 if value is NaN */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator unsigned int() const {
+#else
+  __FP8_HOST__ operator unsigned int() const {
+#endif
     if (internal::hip_fp8_fnuz_is_nan(__x)) {
       return 0;
     }
@@ -1214,7 +1617,11 @@ struct __hip_fp8_e5m2_fnuz {
   }
 
   /*! convert fp8 e5m2 to unsigned long, return 0 if value is NaN */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator unsigned long int() const {
+#else
+  __FP8_HOST__ operator unsigned long int() const {
+#endif
     if (internal::hip_fp8_fnuz_is_nan(__x)) {
       return 0;
     }
@@ -1228,7 +1635,11 @@ struct __hip_fp8_e5m2_fnuz {
   }
 
   /*! convert fp8 e5m2 to unsigned long long, return 0 if value is NaN */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator unsigned long long int() const {
+#else
+  __FP8_HOST__ operator unsigned long long int() const {
+#endif
     if (internal::hip_fp8_fnuz_is_nan(__x)) {
       return 0;
     }
@@ -1242,7 +1653,11 @@ struct __hip_fp8_e5m2_fnuz {
   }
 
   /*! convert fp8 e5m2 to unsigned short, return 0 if value is NaN */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator unsigned short int() const {
+#else
+  __FP8_HOST__ operator unsigned short int() const {
+#endif
     if (internal::hip_fp8_fnuz_is_nan(__x)) {
       return 0;
     }
@@ -1268,31 +1683,59 @@ struct __hip_fp8x2_e5m2_fnuz {
   static constexpr unsigned int __wm = 2;
 
   /*! create fp8x2 e5m2 type from double2 */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8x2_e5m2_fnuz(const double2 val)
+#else
+  __FP8_HOST__ __hip_fp8x2_e5m2_fnuz(const double2 val)
+#endif
       : __x(__hip_cvt_double2_to_fp8x2(val, __default_saturation, __default_interpret)) {}
 
   /*! create fp8x2 e5m2 type from float2 */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8x2_e5m2_fnuz(const float2 val)
+#else
+  __FP8_HOST__ __hip_fp8x2_e5m2_fnuz(const float2 val)
+#endif
       : __x(__hip_cvt_float2_to_fp8x2(val, __default_saturation, __default_interpret)) {}
 
   /*! create fp8x2 e5m2 type from __hip_bfloat162 */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8x2_e5m2_fnuz(const __hip_bfloat162 val)
+#else
+  __FP8_HOST__ __hip_fp8x2_e5m2_fnuz(const __hip_bfloat162 val)
+#endif
       : __x(__hip_cvt_bfloat16raw2_to_fp8x2(val, __default_saturation, __default_interpret)) {}
 
   /*! create fp8x2 e5m2 type from __half2 */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8x2_e5m2_fnuz(const __half2 val)
+#else
+  __FP8_HOST__ __hip_fp8x2_e5m2_fnuz(const __half2 val)
+#endif
       : __x(__hip_cvt_halfraw2_to_fp8x2(val, __default_saturation, __default_interpret)) {}
 
   /*! default construct fp8x2 e5m2 */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8x2_e5m2_fnuz() = default;
+#else
+  __FP8_HOST__ __hip_fp8x2_e5m2_fnuz() = default;
+#endif
 
   /*! convert fp8x2 e5m2 to __half2 */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator __half2() const {
+#else
+  __FP8_HOST__ operator __half2() const {
+#endif
     return __half2(__hip_cvt_fp8x2_to_halfraw2(__x, __default_interpret));
   }
 
   /*! convert fp8x2 e5m2 to float2 */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator float2() const {
+#else
+  __FP8_HOST__ operator float2() const {
+#endif
 #if HIP_FP8_CVT_FAST_PATH
     return internal::cast_to_f32x2_from_f8x2(__x, __default_interpret);
 #else
@@ -1316,7 +1759,11 @@ struct __hip_fp8x4_e5m2_fnuz {
   static constexpr unsigned int __wm = 2;
 
   /*! create fp8x4 e5m2 type from double4 */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8x4_e5m2_fnuz(const double4 val)
+#else
+  __FP8_HOST__ __hip_fp8x4_e5m2_fnuz(const double4 val)
+#endif
       : __x(reinterpret_cast<__hip_fp8x4_storage_t>(
             static_cast<unsigned int>(reinterpret_cast<unsigned char>(__hip_cvt_double_to_fp8(
                                           val.x, __default_saturation, __default_interpret)) |
@@ -1331,7 +1778,11 @@ struct __hip_fp8x4_e5m2_fnuz {
                                           << 24))) {}
 
   /*! create fp8x4 e5m2 type from float4 */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8x4_e5m2_fnuz(const float4 val)
+#else
+  __FP8_HOST__ __hip_fp8x4_e5m2_fnuz(const float4 val)
+#endif
       : __x(reinterpret_cast<__hip_fp8x4_storage_t>(
             static_cast<unsigned int>(reinterpret_cast<unsigned char>(__hip_cvt_float_to_fp8(
                                           val.x, __default_saturation, __default_interpret)) |
@@ -1346,7 +1797,11 @@ struct __hip_fp8x4_e5m2_fnuz {
                                           << 24))) {}
 
   /*! create fp8x4 e5m2 type from two __hip_bfloat162 */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8x4_e5m2_fnuz(const __hip_bfloat162 low, const __hip_bfloat162 high)
+#else
+  __FP8_HOST__ __hip_fp8x4_e5m2_fnuz(const __hip_bfloat162 low, const __hip_bfloat162 high)
+#endif
       : __x(reinterpret_cast<__hip_fp8x4_storage_t>(static_cast<unsigned int>(
             reinterpret_cast<unsigned short>(
                 __hip_cvt_bfloat16raw2_to_fp8x2(high, __default_saturation, __default_interpret)) |
@@ -1355,7 +1810,11 @@ struct __hip_fp8x4_e5m2_fnuz {
                 << 16))) {}
 
   /*! create fp8x4 e5m2 type from two __half2 */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8x4_e5m2_fnuz(const __half2 low, const __half2 high)
+#else
+  __FP8_HOST__ __hip_fp8x4_e5m2_fnuz(const __half2 low, const __half2 high)
+#endif
       : __x(reinterpret_cast<__hip_fp8x4_storage_t>(
             static_cast<unsigned int>(reinterpret_cast<unsigned short>(__hip_cvt_halfraw2_to_fp8x2(
                                           high, __default_saturation, __default_interpret)) |
@@ -1364,10 +1823,18 @@ struct __hip_fp8x4_e5m2_fnuz {
                                           << 16))) {}
 
   /* default construct fp8x4 e5m2 */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ __hip_fp8x4_e5m2_fnuz() = default;
+#else
+  __FP8_HOST__ __hip_fp8x4_e5m2_fnuz() = default;
+#endif
 
   /*! convert fp8x4 e5m2 to float4 */
+#if HIP_FP8_TYPE_FNUZ
   __FP8_HOST_DEVICE__ operator float4() const {
+#else
+  __FP8_HOST__ operator float4() const {
+#endif
     auto x = __x;                                                    // bypass const
     auto fp8x2_low = *reinterpret_cast<__hip_fp8x2_storage_t*>(&x);  // Little E
     auto fp8x2_high = *(reinterpret_cast<__hip_fp8x2_storage_t*>(&x) + 1);
@@ -1388,4 +1855,1079 @@ struct __hip_fp8x4_e5m2_fnuz {
   }
 };
 
-#endif  // _HIP_INCLUDE_HIP_AMD_DETAIL_HIP_FP8_H_
+/**
+ * \brief struct representing ocp fp8 numbers with e4m3 interpretation
+ *
+ * */
+struct __hip_fp8_e4m3 {
+    __hip_fp8_storage_t __x;  //! raw storage of fp8 number
+  constexpr static __hip_saturation_t __default_saturation = __HIP_SATFINITE;
+  constexpr static __hip_fp8_interpretation_t __default_interpret = __HIP_E4M3;
+  constexpr static unsigned int __we = 4;
+  constexpr static unsigned int __wm = 3;
+
+  // TODO: SWDEV-452411
+  // Add cast from unsigned long long, long long to fp8
+
+  /*! create fp8 e4m3 from long */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8_e4m3(const long int val)
+#else
+__FP8_HOST__ __hip_fp8_e4m3(const long int val)
+#endif
+      : __x(__hip_cvt_float_to_fp8(static_cast<float>(val), __default_saturation,
+                                   __default_interpret)) {}
+
+  /*! create fp8 e4m3 from int */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8_e4m3(const int val)
+#else
+__FP8_HOST__ __hip_fp8_e4m3(const int val)
+#endif
+      : __x(__hip_cvt_float_to_fp8(static_cast<float>(val), __default_saturation,
+                                   __default_interpret)) {}
+
+  /*! create fp8 e4m3 from short int */
+  __FP8_HOST_DEVICE__ __hip_fp8_e4m3(const short int val)
+      : __x(__hip_cvt_float_to_fp8(static_cast<float>(val), __default_saturation,
+                                   __default_interpret)) {}
+
+  /*! create fp8 e4m3 from unsigned long */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8_e4m3(const unsigned long int val)
+#else
+__FP8_HOST__ __hip_fp8_e4m3(const unsigned long int val)
+#endif
+      : __x(__hip_cvt_float_to_fp8(static_cast<float>(val), __default_saturation,
+                                   __default_interpret)) {}
+
+  /*! create fp8 e4m3 from unsigned int */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8_e4m3(const unsigned int val)
+#else
+__FP8_HOST__ __hip_fp8_e4m3(const unsigned int val)
+#endif
+      : __x(__hip_cvt_float_to_fp8(static_cast<float>(val), __default_saturation,
+                                   __default_interpret)) {}
+
+  /*! create fp8 e4m3 from unsigned short */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8_e4m3(const unsigned short int val)
+#else
+__FP8_HOST__ __hip_fp8_e4m3(const unsigned short int val)
+#endif
+      : __x(__hip_cvt_float_to_fp8(static_cast<float>(val), __default_saturation,
+                                   __default_interpret)) {}
+
+  /*! create fp8 e4m3 from double */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8_e4m3(const double f)
+#else
+__FP8_HOST__ __hip_fp8_e4m3(const double f)
+#endif
+      : __x(__hip_cvt_double_to_fp8(f, __default_saturation, __default_interpret)) {}
+
+  /*! create fp8 e4m3 from float */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8_e4m3(const float f)
+#else
+__FP8_HOST__ __hip_fp8_e4m3(const float f)
+#endif
+      : __x(__hip_cvt_float_to_fp8(f, __default_saturation, __default_interpret)) {}
+
+  /*! create fp8 e4m3 from __hip_bfloat16 */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8_e4m3(const __hip_bfloat16 f)
+#else
+__FP8_HOST__ __hip_fp8_e4m3(const __hip_bfloat16 f)
+#endif
+      : __x(__hip_cvt_float_to_fp8(static_cast<float>(f), __default_saturation,
+                                   __default_interpret)) {}
+
+  /*! create fp8 e4m3 from __half */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8_e4m3(const __half f)
+#else
+__FP8_HOST__ __hip_fp8_e4m3(const __half f)
+#endif
+      : __x(__hip_cvt_halfraw_to_fp8(static_cast<__half_raw>(f), __default_saturation,
+                                     __default_interpret)) {}
+
+  /*! default construct fp8 e4m3 */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8_e4m3() = default;
+#else
+__FP8_HOST__ __hip_fp8_e4m3() = default;
+#endif
+
+  /*! convert fp8 e4m3 to __half */
+
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator __half() const {
+#else
+__FP8_HOST__ operator __half() const {
+#endif
+    return __half(__hip_cvt_fp8_to_halfraw(__x, __default_interpret));
+  }
+
+  /*! convert fp8 e4m3 to __hip_bfloat16 */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator __hip_bfloat16() const {
+#else
+__FP8_HOST__ operator __hip_bfloat16() const {
+#endif
+    float f = *this;
+    return __hip_bfloat16(f);
+  }
+
+  /*! convert fp8 e4m3 to bool, return false if value is 0, true otherwise */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator bool() const {
+#else
+__FP8_HOST__ operator bool() const {
+#endif
+    // it can be 0x00 (+0.0) since 0x80 will be nan
+    return !(static_cast<unsigned short>(__x) == 0 || static_cast<unsigned short>(__x) == 0x80);
+  }
+
+  /*! convert fp8 e4m3 to char, clamp number to CHAR_MIN/CHAR_MAX if its out of range */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator char() const {
+#else
+__FP8_HOST__ operator char() const {
+#endif
+    if (internal::hip_fp8_ocp_is_nan(__x,__default_interpret)) {
+      return 0;
+    }
+
+    auto fval = internal::cast_from_f8<float, false>(__x, __wm, __we);
+    auto llval = static_cast<long long>(fval);
+    if (llval <= CHAR_MIN) {
+      return CHAR_MIN;
+    } else if (llval >= CHAR_MAX) {
+      return CHAR_MAX;
+    }
+    return static_cast<char>(fval);
+  }
+
+  /*! convert fp8 e4m3 to double */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator double() const {
+#else
+__FP8_HOST__ operator double() const {
+#endif
+    return internal::cast_from_f8<double, false>(__x, __wm, __we);
+  }
+
+  /*! convert fp8 e4m3 to float */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator float() const {
+#else
+__FP8_HOST__ operator float() const {
+#endif
+#if HIP_FP8_CVT_FAST_PATH
+    return internal::cast_to_f32_from_f8(__x, __default_interpret);
+#else
+    return internal::cast_from_f8<float, false>(__x, __wm, __we);
+#endif
+  }
+
+  /*! convert fp8 e4m3 to int, return 0 if value is NaN */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator int() const {
+#else
+__FP8_HOST__ operator int() const {
+#endif
+    if (internal::hip_fp8_ocp_is_nan(__x, __default_interpret)) {
+      return 0;
+    }
+
+    float fval = *this;
+    return static_cast<int>(fval);
+  }
+
+  /*! convert fp8 e4m3 to long, return 0 if value is NaN */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator long int() const {
+#else
+__FP8_HOST__ operator long int() const {
+#endif
+    if (internal::hip_fp8_ocp_is_nan(__x, __default_interpret)) {
+      return 0;
+    }
+
+    float fval = *this;
+    return static_cast<long>(fval);
+  }
+
+  /*! convert fp8 e4m3 to long long, return 0 if value is NaN */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator long long int() const {
+#else
+__FP8_HOST__ operator long long int() const {
+#endif
+    if (internal::hip_fp8_ocp_is_nan(__x, __default_interpret)) {
+      return 0;
+    }
+
+    float fval = *this;
+    return static_cast<long long>(fval);
+  }
+
+  /*! convert fp8 e4m3 to short int, clamp out of bound values, return 0 if value is NaN */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator short int() const {
+#else
+__FP8_HOST__ operator short int() const {
+#endif
+    if (internal::hip_fp8_ocp_is_nan(__x, __default_interpret)) {
+      return 0;
+    }
+
+    float fval = *this;
+    auto llval = static_cast<long long>(fval);
+    if (llval <= SHRT_MIN) {
+      return SHRT_MIN;
+    } else if (llval >= SHRT_MAX) {
+      return SHRT_MAX;
+    }
+    return static_cast<short>(fval);
+  }
+
+  /*! convert fp8 e4m3 to signed char, clamp out of bound values, return 0 if value is NaN */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator signed char() const {
+#else
+__FP8_HOST__ operator signed char() const {
+#endif
+    if (internal::hip_fp8_ocp_is_nan(__x, __default_interpret)) {
+      return 0;
+    }
+
+    float fval = *this;
+    auto llval = static_cast<long long>(fval);
+    if (llval <= SCHAR_MIN) {
+      return SCHAR_MIN;
+    } else if (llval >= SCHAR_MAX) {
+      return SCHAR_MAX;
+    }
+    return static_cast<signed char>(fval);
+  }
+
+  /*! convert fp8 e4m3 to unsigned char, clamp out of bound values, return 0 if value is NaN */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator unsigned char() const {
+#else
+__FP8_HOST__ operator unsigned char() const {
+#endif
+    if (internal::hip_fp8_ocp_is_nan(__x, __default_interpret)) {
+      return 0;
+    }
+
+    float fval = *this;
+    auto llval = static_cast<long long>(fval);
+    if (llval <= 0) {
+      return 0;
+    } else if (llval >= UCHAR_MAX) {
+      return UCHAR_MAX;
+    }
+    return static_cast<unsigned char>(fval);
+  }
+
+  /*! convert fp8 e4m3 to unsigned int, return 0 if value is NaN */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator unsigned int() const {
+#else
+__FP8_HOST__ operator unsigned int() const {
+#endif
+    if (internal::hip_fp8_ocp_is_nan(__x, __default_interpret)) {
+      return 0;
+    }
+
+    float fval = *this;
+    auto llval = static_cast<long long>(fval);
+    if (llval <= 0) {
+      return 0;
+    }
+    return static_cast<unsigned int>(fval);
+  }
+
+  /*! convert fp8 e4m3 to unsigned long, return 0 if value is NaN */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator unsigned long int() const {
+#else
+__FP8_HOST__ operator unsigned long int() const {
+#endif
+    if (internal::hip_fp8_ocp_is_nan(__x, __default_interpret)) {
+      return 0;
+    }
+
+    float fval = *this;
+    auto llval = static_cast<long long>(fval);
+    if (llval <= 0) {
+      return 0;
+    }
+    return static_cast<unsigned long>(fval);
+  }
+
+  /*! convert fp8 e4m3 to long long int, return 0 if value is NaN */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator unsigned long long int() const {
+#else
+__FP8_HOST__ operator unsigned long long int() const {
+#endif
+    if (internal::hip_fp8_ocp_is_nan(__x, __default_interpret)) {
+      return 0;
+    }
+
+    float fval = *this;
+    auto llval = static_cast<long long>(fval);
+    if (llval <= 0) {
+      return 0;
+    }
+    return static_cast<unsigned long long>(fval);
+  }
+
+  /*! convert fp8 e4m3 to unsigned short, return 0 if value is NaN */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator unsigned short int() const {
+#else
+__FP8_HOST__ operator unsigned short int() const {
+#endif
+    if (internal::hip_fp8_ocp_is_nan(__x,__default_interpret)) {
+      return 0;
+    }
+
+    float fval = *this;
+    auto llval = static_cast<long long>(fval);
+    if (llval <= 0) {
+      return 0;
+    }
+    return static_cast<unsigned short>(fval);
+  }
+};
+
+/**
+ * \brief struct representing two ocp fp8 numbers with e4m3 interpretation
+ *
+ * */
+struct __hip_fp8x2_e4m3 {
+  __hip_fp8x2_storage_t __x;  //! raw storage of two fp8 numbers
+  static constexpr __hip_saturation_t __default_saturation = __HIP_SATFINITE;
+  static constexpr __hip_fp8_interpretation_t __default_interpret = __HIP_E4M3;
+  static constexpr unsigned int __we = 4;
+  static constexpr unsigned int __wm = 3;
+
+  /*! create fp8x2 e4m3 type from double2 */
+
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8x2_e4m3(const double2 val)
+#else
+__FP8_HOST__ __hip_fp8x2_e4m3(const double2 val)
+#endif
+      : __x(__hip_cvt_double2_to_fp8x2(val, __default_saturation, __default_interpret)) {}
+
+  /*! create fp8x2 e4m3 type from float2 */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8x2_e4m3(const float2 val)
+#else
+__FP8_HOST__ __hip_fp8x2_e4m3(const float2 val)
+#endif
+      : __x(__hip_cvt_float2_to_fp8x2(val, __default_saturation, __default_interpret)) {}
+
+  /*! create fp8x2 e4m3 type from __hip_bfloat162 */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8x2_e4m3(const __hip_bfloat162 val)
+#else
+__FP8_HOST__ __hip_fp8x2_e4m3(const __hip_bfloat162 val)
+#endif
+      : __x(__hip_cvt_bfloat16raw2_to_fp8x2(val, __default_saturation, __default_interpret)) {}
+
+  /*! create fp8x2 e4m3 type from __half2 */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8x2_e4m3(const __half2 val)
+#else
+__FP8_HOST__ __hip_fp8x2_e4m3(const __half2 val)
+#endif
+      : __x(__hip_cvt_halfraw2_to_fp8x2(val, __default_saturation, __default_interpret)) {}
+
+  /*! Default construct of fp8x2 e4m3 */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8x2_e4m3() = default;
+#else
+__FP8_HOST__ __hip_fp8x2_e4m3() = default;
+#endif
+
+  /*! convert fp8x2 e4m3 to __half2 */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator __half2() const {
+#else
+__FP8_HOST__ operator __half2() const {
+#endif
+    return __half2(__hip_cvt_fp8x2_to_halfraw2(__x, __default_interpret));
+  }
+
+  /*! convert fp8x2 e4m3 to float2 */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator float2() const {
+#else
+__FP8_HOST__ operator float2() const {
+#endif
+#if HIP_FP8_CVT_FAST_PATH
+    return internal::cast_to_f32x2_from_f8x2(__x, __default_interpret);
+#else
+    return float2(internal::cast_from_f8<float, false>(static_cast<__hip_fp8_storage_t>(__x & 0xFF), __wm, __we),
+                  internal::cast_from_f8<float, false>(static_cast<__hip_fp8_storage_t>(__x >> 8), __wm, __we));
+#endif
+  }
+};
+
+/**
+ * \brief struct representing four ocp fp8 numbers with e4m3 interpretation
+ *
+ * */
+struct __hip_fp8x4_e4m3 {
+  __hip_fp8x4_storage_t __x;  //! raw storage of four fp8 numbers
+  static constexpr __hip_saturation_t __default_saturation = __HIP_SATFINITE;
+  static constexpr __hip_fp8_interpretation_t __default_interpret = __HIP_E4M3;
+  static constexpr unsigned int __we = 4;
+  static constexpr unsigned int __wm = 3;
+
+  /*! create fp8x4 e4m3 type from double4 */
+
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8x4_e4m3(const double4 val)
+#else
+__FP8_HOST__ __hip_fp8x4_e4m3(const double4 val)
+#endif
+      : __x{reinterpret_cast<__hip_fp8x4_storage_t>(
+            static_cast<unsigned int>(reinterpret_cast<unsigned char>(__hip_cvt_double_to_fp8(
+                                          val.x, __default_saturation, __default_interpret)) |
+                                      reinterpret_cast<unsigned char>(__hip_cvt_double_to_fp8(
+                                          val.y, __default_saturation, __default_interpret))
+                                          << 8 |
+                                      reinterpret_cast<unsigned char>(__hip_cvt_double_to_fp8(
+                                          val.z, __default_saturation, __default_interpret))
+                                          << 16 |
+                                      reinterpret_cast<unsigned char>(__hip_cvt_double_to_fp8(
+                                          val.w, __default_saturation, __default_interpret))
+                                          << 24))} {}
+
+  /*! create fp8x4 e4m3 type from float4 */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8x4_e4m3(const float4 val)
+#else
+__FP8_HOST__ __hip_fp8x4_e4m3(const float4 val)
+#endif
+      : __x{reinterpret_cast<__hip_fp8x4_storage_t>(
+            static_cast<unsigned int>(reinterpret_cast<unsigned char>(__hip_cvt_float_to_fp8(
+                                          val.x, __default_saturation, __default_interpret)) |
+                                      reinterpret_cast<unsigned char>(__hip_cvt_float_to_fp8(
+                                          val.y, __default_saturation, __default_interpret))
+                                          << 8 |
+                                      reinterpret_cast<unsigned char>(__hip_cvt_float_to_fp8(
+                                          val.z, __default_saturation, __default_interpret))
+                                          << 16 |
+                                      reinterpret_cast<unsigned char>(__hip_cvt_float_to_fp8(
+                                          val.w, __default_saturation, __default_interpret))
+                                          << 24))} {}
+
+  /*! create fp8x4 e4m3 type from two __hip_bfloat162 */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8x4_e4m3(const __hip_bfloat162 low, const __hip_bfloat162 high)
+#else
+__FP8_HOST__ __hip_fp8x4_e4m3(const __hip_bfloat162 low, const __hip_bfloat162 high)
+#endif
+      : __x(reinterpret_cast<__hip_fp8x4_storage_t>(static_cast<unsigned int>(
+            reinterpret_cast<unsigned short>(
+                __hip_cvt_bfloat16raw2_to_fp8x2(high, __default_saturation, __default_interpret)) |
+            reinterpret_cast<unsigned short>(
+                __hip_cvt_bfloat16raw2_to_fp8x2(low, __default_saturation, __default_interpret))
+                << 16))) {}
+
+  /*! create fp8x4 e4m3 type from two __half2 */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8x4_e4m3(const __half2 low, const __half2 high)
+#else
+__FP8_HOST__ __hip_fp8x4_e4m3(const __half2 low, const __half2 high)
+#endif
+      : __x(reinterpret_cast<__hip_fp8x4_storage_t>(
+            static_cast<unsigned int>(reinterpret_cast<unsigned short>(__hip_cvt_halfraw2_to_fp8x2(
+                                          high, __default_saturation, __default_interpret)) |
+                                      reinterpret_cast<unsigned short>(__hip_cvt_halfraw2_to_fp8x2(
+                                          low, __default_saturation, __default_interpret))
+                                          << 16))) {}
+
+  /*! Default construct fp8x4 e4m3 */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8x4_e4m3() = default;
+#else
+__FP8_HOST__ __hip_fp8x4_e4m3() = default;
+#endif
+
+  /*! convert fp8x4 e4m3 to float4 */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator float4() const {
+#else
+__FP8_HOST__ operator float4() const {
+#endif
+    auto x = __x;                                                    // bypass const
+    auto fp8x2_low = *reinterpret_cast<__hip_fp8x2_storage_t*>(&x);  // Little E
+    auto fp8x2_high = *(reinterpret_cast<__hip_fp8x2_storage_t*>(&x) + 1);
+#if HIP_FP8_CVT_FAST_PATH
+    float2 high = internal::cast_to_f32x2_from_f8x2(fp8x2_high, __default_interpret);
+    float2 low = internal::cast_to_f32x2_from_f8x2(fp8x2_low, __default_interpret);
+#else
+    float2 high = float2(internal::cast_from_f8<float, false>(
+                             static_cast<__hip_fp8_storage_t>((fp8x2_high << 8) >> 8), __wm, __we),
+                         internal::cast_from_f8<float, false>(
+                             static_cast<__hip_fp8_storage_t>(fp8x2_high >> 8), __wm, __we));
+    float2 low = float2(internal::cast_from_f8<float, false>(
+                            static_cast<__hip_fp8_storage_t>((fp8x2_low << 8) >> 8), __wm, __we),
+                        internal::cast_from_f8<float, false>(
+                            static_cast<__hip_fp8_storage_t>(fp8x2_low >> 8), __wm, __we));
+#endif
+    return float4(low.x, low.y, high.x, high.y);
+  }
+};
+
+/**
+ * \brief struct representing  ocp fp8 numbers with e5m2 interpretation
+ *
+ * */
+struct __hip_fp8_e5m2 {
+  __hip_fp8_storage_t __x;  //! raw storage of one fp8 numbers
+  static constexpr __hip_saturation_t __default_saturation = __HIP_SATFINITE;
+  static constexpr __hip_fp8_interpretation_t __default_interpret = __HIP_E5M2;
+  static constexpr unsigned int __we = 5;
+  static constexpr unsigned int __wm = 2;
+
+
+  // TODO: SWDEV-452411
+  // Add cast from unsigned long long, long long to fp8
+
+  /*! create fp8 e5m2 type from long */
+
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8_e5m2(const long int val)
+#else
+__FP8_HOST__ __hip_fp8_e5m2(const long int val)
+#endif
+      : __x(__hip_cvt_float_to_fp8(static_cast<float>(val), __default_saturation,
+                                   __default_interpret)) {}
+
+  /*! create fp8 e5m2 type from int */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8_e5m2(const int val)
+#else
+__FP8_HOST__ __hip_fp8_e5m2(const int val)
+#endif
+      : __x(__hip_cvt_float_to_fp8(static_cast<float>(val), __default_saturation,
+                                   __default_interpret)) {}
+
+  /*! create fp8 e5m2 type from short int */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8_e5m2(const short int val)
+#else
+__FP8_HOST__ __hip_fp8_e5m2(const short int val)
+#endif
+      : __x(__hip_cvt_float_to_fp8(static_cast<float>(val), __default_saturation,
+                                   __default_interpret)) {}
+
+  /*! create fp8 e5m2 type from unsigned long */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8_e5m2(const unsigned long int val)
+#else
+__FP8_HOST__ __hip_fp8_e5m2(const unsigned long int val)
+#endif
+      : __x(__hip_cvt_float_to_fp8(static_cast<float>(val), __default_saturation,
+                                   __default_interpret)) {}
+
+  /*! create fp8 e5m2 type from unsigned int */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8_e5m2(const unsigned int val)
+#else
+__FP8_HOST__ __hip_fp8_e5m2(const unsigned int val)
+#endif
+      : __x(__hip_cvt_float_to_fp8(static_cast<float>(val), __default_saturation,
+                                   __default_interpret)) {}
+
+  /*! create fp8 e5m2 type from unsigned short */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8_e5m2(const unsigned short int val)
+#else
+__FP8_HOST__ __hip_fp8_e5m2(const unsigned short int val)
+#endif
+      : __x(__hip_cvt_float_to_fp8(static_cast<float>(val), __default_saturation,
+                                   __default_interpret)) {}
+
+  /*! create fp8 e5m2 type from double */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8_e5m2(const double f)
+#else
+__FP8_HOST__ __hip_fp8_e5m2(const double f)
+#endif
+      : __x(__hip_cvt_double_to_fp8(f, __default_saturation, __default_interpret)) {}
+
+  /*! create fp8 e5m2 type from float */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8_e5m2(const float f)
+#else
+__FP8_HOST__ __hip_fp8_e5m2(const float f)
+#endif
+      : __x(__hip_cvt_float_to_fp8(f, __default_saturation, __default_interpret)) {}
+
+  /*! create fp8 e5m2 type from __hip_bfloat16 */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8_e5m2(const __hip_bfloat16 f)
+#else
+__FP8_HOST__ __hip_fp8_e5m2(const __hip_bfloat16 f)
+#endif
+      : __x(__hip_cvt_float_to_fp8(static_cast<float>(f), __default_saturation,
+                                   __default_interpret)) {}
+
+  /*! create fp8 e5m2 type from __hip_bfloat16 */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8_e5m2(const __half f)
+#else
+__FP8_HOST__ __hip_fp8_e5m2(const __half f)
+#endif
+      : __x(__hip_cvt_halfraw_to_fp8(static_cast<__half_raw>(f), __default_saturation,
+                                     __default_interpret)) {}
+
+  /*! default construct fp8 e5m2 */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8_e5m2() = default;
+#else
+__FP8_HOST__ __hip_fp8_e5m2() = default;
+#endif
+
+  /*! convert fp8 e5m2 to float */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator float() const {
+#else
+__FP8_HOST__ operator float() const {
+#endif
+#if HIP_FP8_CVT_FAST_PATH
+    return internal::cast_to_f32_from_f8(__x, __default_interpret);
+#else
+    return internal::cast_from_f8<float, false>(__x, __wm, __we, __default_saturation == __HIP_SATFINITE);
+#endif
+  }
+
+  /*! convert fp8 e5m2 to __half */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator __half() const {
+#else
+__FP8_HOST__ operator __half() const {
+#endif
+    return __half(__hip_cvt_fp8_to_halfraw(__x, __default_interpret));
+  }
+
+  /*! convert fp8 e5m2 to __hip_bfloat16 */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator __hip_bfloat16() const {
+#else
+__FP8_HOST__ operator __hip_bfloat16() const {
+#endif
+    float f = *this;
+    return __hip_bfloat16(f);
+  }
+
+  /*! convert fp8 e4m3 to bool, return false if value is 0, true otherwise */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator bool() const {
+#else
+__FP8_HOST__ operator bool() const {
+#endif
+    // it can be 0x00 (+0.0) since 0x80 will be nan
+    return !(static_cast<unsigned short>(__x) == 0 || static_cast<unsigned short>(__x) == 0x80);
+  }
+
+  /*! convert fp8 e5m2 to char, clamp out of bound values, return 0 if value is NaN */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator char() const {
+#else
+__FP8_HOST__ operator char() const {
+#endif
+    if (internal::hip_fp8_ocp_is_nan(__x, __default_interpret)) {
+      return 0;
+    }
+
+    float fval = *this;
+    auto llval = static_cast<long long>(fval);
+    if (llval <= CHAR_MIN) {
+      return CHAR_MIN;
+    } else if (llval >= CHAR_MAX) {
+      return CHAR_MAX;
+    }
+    return static_cast<char>(fval);
+  }
+
+  /*! convert fp8 e5m2 to double */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator double() const {
+#else
+__FP8_HOST__ operator double() const {
+#endif
+    return internal::cast_from_f8<double, false>(__x, __wm, __we, __default_saturation == __HIP_SATFINITE);
+  }
+
+  /*! convert fp8 e5m2 to int, return 0 if value is NaN */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator int() const {
+#else
+__FP8_HOST__ operator int() const {
+#endif
+    if (internal::hip_fp8_ocp_is_nan(__x, __default_interpret)) {
+      return 0;
+    }
+
+    float fval = *this;
+    return static_cast<int>(fval);
+  }
+
+  /*! convert fp8 e5m2 to long, return 0 if value is NaN */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator long int() const {
+#else
+__FP8_HOST__ operator long int() const {
+#endif
+    if (internal::hip_fp8_ocp_is_nan(__x, __default_interpret)) {
+      return 0;
+    }
+
+    float fval = *this;
+    return static_cast<long>(fval);
+  }
+
+  /*! convert fp8 e5m2 to long long, return 0 if value is NaN */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator long long int() const {
+#else
+__FP8_HOST__ operator long long int() const {
+#endif
+    if (internal::hip_fp8_ocp_is_nan(__x, __default_interpret)) {
+      return 0;
+    }
+
+    float fval = *this;
+    return static_cast<long long>(fval);
+  }
+
+  /*! convert fp8 e5m2 to short, clamp out of bound values, return 0 if value is NaN */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator short int() const {
+#else
+__FP8_HOST__ operator short int() const {
+#endif
+    if (internal::hip_fp8_ocp_is_nan(__x, __default_interpret)) {
+      return 0;
+    }
+
+    float fval = *this;
+    auto llval = static_cast<long long>(fval);
+    if (llval <= SHRT_MIN) {
+      return SHRT_MIN;
+    } else if (llval >= SHRT_MAX) {
+      return SHRT_MAX;
+    }
+    return static_cast<short>(fval);
+  }
+
+  /*! convert fp8 e5m2 to signed char, clamp out of bound values, return 0 if value is NaN */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator signed char() const {
+#else
+__FP8_HOST__ operator signed char() const {
+#endif
+    if (internal::hip_fp8_ocp_is_nan(__x, __default_interpret)) {
+      return 0;
+    }
+
+    float fval = *this;
+    auto llval = static_cast<long long>(fval);
+    if (llval <= SCHAR_MIN) {
+      return SCHAR_MIN;
+    } else if (llval >= SCHAR_MAX) {
+      return SCHAR_MAX;
+    }
+    return static_cast<signed char>(fval);
+  }
+
+  /*! convert fp8 e5m2 to unsigned char, clamp out of bound values, return 0 if value is NaN */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator unsigned char() const {
+#else
+__FP8_HOST__ operator unsigned char() const {
+#endif
+    if (internal::hip_fp8_ocp_is_nan(__x, __default_interpret)) {
+      return 0;
+    }
+
+    float fval = *this;
+    auto llval = static_cast<long long>(fval);
+    if (llval <= 0) {
+      return 0;
+    } else if (llval >= UCHAR_MAX) {
+      return UCHAR_MAX;
+    }
+    return static_cast<unsigned char>(fval);
+  }
+
+  /*! convert fp8 e5m2 to unsigned int, return 0 if value is NaN */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator unsigned int() const {
+#else
+__FP8_HOST__ operator unsigned int() const {
+#endif
+    if (internal::hip_fp8_ocp_is_nan(__x, __default_interpret)) {
+      return 0;
+    }
+
+    float fval = *this;
+    auto llval = static_cast<long long>(fval);
+    if (llval <= 0) {
+      return 0;
+    }
+    return static_cast<unsigned int>(fval);
+  }
+
+  /*! convert fp8 e5m2 to unsigned long, return 0 if value is NaN */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator unsigned long int() const {
+#else
+__FP8_HOST__ operator unsigned long int() const {
+#endif
+    if (internal::hip_fp8_ocp_is_nan(__x, __default_interpret)) {
+      return 0;
+    }
+
+    float fval = *this;
+    auto llval = static_cast<long long>(fval);
+    if (llval <= 0) {
+      return 0;
+    }
+    return static_cast<unsigned long>(fval);
+  }
+
+  /*! convert fp8 e5m2 to unsigned long long, return 0 if value is NaN */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator unsigned long long int() const {
+#else
+__FP8_HOST__ operator unsigned long long int() const {
+#endif
+    if (internal::hip_fp8_ocp_is_nan(__x, __default_interpret)) {
+      return 0;
+    }
+
+    float fval = *this;
+    auto llval = static_cast<long long>(fval);
+    if (llval <= 0) {
+      return 0;
+    }
+    return static_cast<unsigned long long>(fval);
+  }
+
+  /*! convert fp8 e5m2 to unsigned short, return 0 if value is NaN */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator unsigned short int() const {
+#else
+__FP8_HOST__ operator unsigned short int() const {
+#endif
+    if (internal::hip_fp8_ocp_is_nan(__x, __default_interpret)) {
+      return 0;
+    }
+
+    float fval = *this;
+    auto llval = static_cast<long long>(fval);
+    if (llval <= 0) {
+      return 0;
+    }
+    return static_cast<unsigned short>(fval);
+  }
+};
+
+/**
+ * \brief struct representing two ocp fp8 numbers with e5m2 interpretation
+ *
+ */
+struct __hip_fp8x2_e5m2 {
+  __hip_fp8x2_storage_t __x;  //! raw storage of two fp8 numbers
+  static constexpr __hip_saturation_t __default_saturation = __HIP_SATFINITE;
+  static constexpr __hip_fp8_interpretation_t __default_interpret = __HIP_E5M2;
+  static constexpr unsigned int __we = 5;
+  static constexpr unsigned int __wm = 2;
+
+  /*! create fp8x2 e5m2 type from double2 */
+
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8x2_e5m2(const double2 val)
+#else
+__FP8_HOST__ __hip_fp8x2_e5m2(const double2 val)
+#endif
+      : __x(__hip_cvt_double2_to_fp8x2(val, __default_saturation, __default_interpret)) {}
+
+  /*! create fp8x2 e5m2 type from float2 */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8x2_e5m2(const float2 val)
+#else
+__FP8_HOST__ __hip_fp8x2_e5m2(const float2 val)
+#endif
+      : __x(__hip_cvt_float2_to_fp8x2(val, __default_saturation, __default_interpret)) {}
+
+  /*! create fp8x2 e5m2 type from __hip_bfloat162 */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8x2_e5m2(const __hip_bfloat162 val)
+#else
+__FP8_HOST__ __hip_fp8x2_e5m2(const __hip_bfloat162 val)
+#endif
+      : __x(__hip_cvt_bfloat16raw2_to_fp8x2(val, __default_saturation, __default_interpret)) {}
+
+  /*! create fp8x2 e5m2 type from __half2 */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8x2_e5m2(const __half2 val)
+#else
+__FP8_HOST__ __hip_fp8x2_e5m2(const __half2 val)
+#endif
+      : __x(__hip_cvt_halfraw2_to_fp8x2(val, __default_saturation, __default_interpret)) {}
+
+  /*! default construct fp8x2 e5m2 */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8x2_e5m2() = default;
+#else
+__FP8_HOST__ __hip_fp8x2_e5m2() = default;
+#endif
+
+  /*! convert fp8x2 e5m2 to __half2 */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator __half2() const {
+#else
+__FP8_HOST__ operator __half2() const {
+#endif
+    return __half2(__hip_cvt_fp8x2_to_halfraw2(__x, __default_interpret));
+  }
+
+  /*! convert fp8x2 e5m2 to float2 */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator float2() const {
+#else
+__FP8_HOST__ operator float2() const {
+#endif
+#if HIP_FP8_CVT_FAST_PATH
+    return internal::cast_to_f32x2_from_f8x2(__x, __default_interpret);
+#else
+    return float2(internal::cast_from_f8<float, false>(static_cast<__hip_fp8_storage_t>(__x & 0xFF), __wm, __we, __default_saturation == __HIP_SATFINITE),
+                  internal::cast_from_f8<float, false>(static_cast<__hip_fp8_storage_t>(__x >> 8), __wm, __we, __default_saturation == __HIP_SATFINITE));
+#endif
+  }
+};
+
+/**
+ * \brief struct representing four ocp fp8 numbers with e5m2 interpretation
+ *
+ */
+struct __hip_fp8x4_e5m2 {
+  __hip_fp8x4_storage_t __x;  //! raw storage of four fp8 numbers
+  static constexpr __hip_saturation_t __default_saturation = __HIP_SATFINITE;
+  static constexpr __hip_fp8_interpretation_t __default_interpret = __HIP_E5M2;
+  static constexpr unsigned int __we = 5;
+  static constexpr unsigned int __wm = 2;
+
+  /*! create fp8x4 e5m2 type from double4 */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8x4_e5m2(const double4 val)
+#else
+__FP8_HOST__ __hip_fp8x4_e5m2(const double4 val)
+#endif
+      : __x(reinterpret_cast<__hip_fp8x4_storage_t>(
+            static_cast<unsigned int>(reinterpret_cast<unsigned char>(__hip_cvt_double_to_fp8(
+                                          val.x, __default_saturation, __default_interpret)) |
+                                      reinterpret_cast<unsigned char>(__hip_cvt_double_to_fp8(
+                                          val.y, __default_saturation, __default_interpret))
+                                          << 8 |
+                                      reinterpret_cast<unsigned char>(__hip_cvt_double_to_fp8(
+                                          val.z, __default_saturation, __default_interpret))
+                                          << 16 |
+                                      reinterpret_cast<unsigned char>(__hip_cvt_double_to_fp8(
+                                          val.w, __default_saturation, __default_interpret))
+                                          << 24))) {}
+
+  /*! create fp8x4 e5m2 type from float4 */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8x4_e5m2(const float4 val)
+#else
+__FP8_HOST__ __hip_fp8x4_e5m2(const float4 val)
+#endif
+      : __x(reinterpret_cast<__hip_fp8x4_storage_t>(
+            static_cast<unsigned int>(reinterpret_cast<unsigned char>(__hip_cvt_float_to_fp8(
+                                          val.x, __default_saturation, __default_interpret)) |
+                                      reinterpret_cast<unsigned char>(__hip_cvt_float_to_fp8(
+                                          val.y, __default_saturation, __default_interpret))
+                                          << 8 |
+                                      reinterpret_cast<unsigned char>(__hip_cvt_float_to_fp8(
+                                          val.z, __default_saturation, __default_interpret))
+                                          << 16 |
+                                      reinterpret_cast<unsigned char>(__hip_cvt_float_to_fp8(
+                                          val.w, __default_saturation, __default_interpret))
+                                          << 24))) {}
+
+  /*! create fp8x4 e5m2 type from two __hip_bfloat162 */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8x4_e5m2(const __hip_bfloat162 low, const __hip_bfloat162 high)
+#else
+__FP8_HOST__ __hip_fp8x4_e5m2(const __hip_bfloat162 low, const __hip_bfloat162 high)
+#endif
+      : __x(reinterpret_cast<__hip_fp8x4_storage_t>(static_cast<unsigned int>(
+            reinterpret_cast<unsigned short>(
+                __hip_cvt_bfloat16raw2_to_fp8x2(high, __default_saturation, __default_interpret)) |
+            reinterpret_cast<unsigned short>(
+                __hip_cvt_bfloat16raw2_to_fp8x2(low, __default_saturation, __default_interpret))
+                << 16))) {}
+
+  /*! create fp8x4 e5m2 type from two __half2 */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8x4_e5m2(const __half2 low, const __half2 high)
+#else
+__FP8_HOST__ __hip_fp8x4_e5m2(const __half2 low, const __half2 high)
+#endif
+      : __x(reinterpret_cast<__hip_fp8x4_storage_t>(
+            static_cast<unsigned int>(reinterpret_cast<unsigned short>(__hip_cvt_halfraw2_to_fp8x2(
+                                          high, __default_saturation, __default_interpret)) |
+                                      reinterpret_cast<unsigned short>(__hip_cvt_halfraw2_to_fp8x2(
+                                          low, __default_saturation, __default_interpret))
+                                          << 16))) {}
+
+  /* default construct fp8x4 e5m2 */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ __hip_fp8x4_e5m2() = default;
+#else
+__FP8_HOST__ __hip_fp8x4_e5m2() = default;
+#endif
+
+  /*! convert fp8x4 e5m2 to float4 */
+#if HIP_FP8_TYPE_OCP
+__FP8_HOST_DEVICE__ operator float4() const {
+#else
+__FP8_HOST__ operator float4() const {
+#endif
+    auto x = __x;                                                    // bypass const
+    auto fp8x2_low = *reinterpret_cast<__hip_fp8x2_storage_t*>(&x);  // Little E
+    auto fp8x2_high = *(reinterpret_cast<__hip_fp8x2_storage_t*>(&x) + 1);
+#if HIP_FP8_CVT_FAST_PATH
+    float2 high = internal::cast_to_f32x2_from_f8x2(fp8x2_high, __default_interpret);
+    float2 low = internal::cast_to_f32x2_from_f8x2(fp8x2_low, __default_interpret);
+#else
+    float2 high = float2(internal::cast_from_f8<float, false>(
+                             static_cast<__hip_fp8_storage_t>((fp8x2_high << 8) >> 8), __wm, __we, __default_saturation == __HIP_SATFINITE),
+                         internal::cast_from_f8<float, false>(
+                             static_cast<__hip_fp8_storage_t>(fp8x2_high >> 8), __wm, __we, __default_saturation == __HIP_SATFINITE));
+    float2 low = float2(internal::cast_from_f8<float, false>(
+                            static_cast<__hip_fp8_storage_t>((fp8x2_low << 8) >> 8), __wm, __we,  __default_saturation == __HIP_SATFINITE),
+                        internal::cast_from_f8<float, false>(
+                            static_cast<__hip_fp8_storage_t>(fp8x2_low >> 8), __wm, __we, __default_saturation == __HIP_SATFINITE));
+#endif
+    return float4(low.x, low.y, high.x, high.y);
+  }
+};
+#endif  // _HIP_INCLUDE_HIP_AMD_DETAIL_HIP_FP8_H_
\ No newline at end of file