yassinsolim · Nov 4, 2026 · Oct 27, 2026 · Oct 28, 2026 · Oct 28, 2026 · Oct 29, 2026
@@ -21,6 +21,10 @@ Full documentation for MIOpen is available [here](https://rocm.docs.amd.com/proj
   `DeviceGroupedConvBwdWeight_Wmma_CShuffleV3<96, 96, 96, 48, ...>` entry the system database
   served for it. That kernel returned a weight gradient with a relative error of about 0.306.
   The problem now uses the next recorded solver, `GemmWrwUniversal`.
+* [FP8] Fixed wrong FP8 and BF8 results on gfx942 from a library built for several targets that
+  include gfx950 or gfx12 ([#14375](https://github.com/ROCm/rocm-libraries/issues/14375)). The
+  FP8 exponent bias was picked once at build time from `GPU_TARGETS`. MIOpen now picks it at run
+  time from the device: the IEEE bias on gfx950 and gfx12, the FNUZ bias everywhere else.
 
 ### Added
 * [Conv] Added `ConvDirectDepthwiseFwd2D`, a direct 2D depthwise forward convolution solver for
@@ -35,6 +39,8 @@ Full documentation for MIOpen is available [here](https://rocm.docs.amd.com/proj
 ### Removed
 * Removed the unused `miopenOperatorDescriptor` type from `miopen.h` (source-level only, ABI-neutral).
 * Removed `miopenAlphaBetaCase_t` and its unprefixed enumerators `DEFAULT`, `SCALE`, `BILINEAR` and `ERROR_STATE` from `miopen.h` (source-level only, ABI-neutral).
+* Removed the `MIOPEN_FP8_IEEE_EXPONENT_BIAS` CMake option and its `config.h` define. The FP8
+  exponent bias now follows the device at run time.
 
 ## MIOpen 3.6.2 for ROCm 10.2.0
 

@@ -22,12 +22,6 @@ function(_miopen_arch_matches arch_value target_list out_var)
     set(${out_var} FALSE PARENT_SCOPE)
 endfunction()
 
-# Returns TRUE in out_var if GPU_TARGETS matches any entry in the given target list.
-function(_miopen_gpu_targets_match target_list out_var)
-    _miopen_arch_matches("${GPU_TARGETS}" ${target_list} _matched)
-    set(${out_var} ${_matched} PARENT_SCOPE)
-endfunction()
-
 # Apply CK compile definitions to a target so its CK macro set matches how CK
 # itself was built (see composable_kernel/CMakeLists.txt).
 #
@@ -63,15 +57,3 @@ function(miopen_apply_ck_compile_definitions target)
         target_compile_definitions(${target} PRIVATE CK_USE_OCP_FP8)
     endif()
 endfunction()
-
-# composable_kernel OCP FP8 mode requires IEEE FP8 exponent bias; keep MIOpen's
-# FP8 semantics aligned by defaulting MIOPEN_FP8_IEEE_EXPONENT_BIAS to ON when
-# needed. This file is included before src/CMakeLists.txt declares option(), so
-# seeding the cache here (without FORCE) both takes effect and still lets
-# advanced users override via -DMIOPEN_FP8_IEEE_EXPONENT_BIAS=...
-if(MIOPEN_USE_COMPOSABLEKERNEL)
-    _miopen_gpu_targets_match(MIOPEN_CK_USE_OCP_FP8_TARGETS _miopen_ck_use_ocp_fp8)
-    if(_miopen_ck_use_ocp_fp8)
-        set(MIOPEN_FP8_IEEE_EXPONENT_BIAS ON CACHE BOOL "Sets the FP8 exponent bias to IEEE")
-    endif()
-endif()
@@ -551,11 +551,15 @@ template <>
 inline void Driver::InitDataType<float8_fnuz>()
 {
     data_type = miopenFloat8_fnuz;
+    miopen_f8::set_hip_f8_ieee_exponent_bias(
+        miopen::UsesIeeeExponentBias(miopen::deref(GetHandle()).GetDeviceName()));
 }
 template <>
 inline void Driver::InitDataType<bfloat8_fnuz>()
 {
     data_type = miopenBFloat8_fnuz;
+    miopen_f8::set_hip_f8_ieee_exponent_bias(
+        miopen::UsesIeeeExponentBias(miopen::deref(GetHandle()).GetDeviceName()));
 }
 // "std::is_same<Tgpu, float>{}" used to avoid "static_assert" compilation error,
 // which occurs when the condition does not depend in any way on the template parameters.

@@ -54,7 +54,6 @@
 //     remain in the future)  perform final conversion (and rounding) of FP32
 //     to BF16 results. This affects the main functionality of the library.
 #cmakedefine01 MIOPEN_USE_RNE_BFLOAT16
-#cmakedefine01 MIOPEN_FP8_IEEE_EXPONENT_BIAS
 #cmakedefine01 MIOPEN_FP8_CLIPPING
 
 // clang-format on

@@ -22,12 +22,8 @@ endif()
 #     remain in the future)  perform final conversion (and rounding) of FP32
 #     to BF16 results. This affects the main functionality of the library.
 option( MIOPEN_USE_RNE_BFLOAT16 "Sets rounding scheme for bfloat16 type" ON )
-option( MIOPEN_FP8_IEEE_EXPONENT_BIAS "Sets the FP8 exponent bias to IEEE" OFF)
 option( MIOPEN_FP8_CLIPPING "Sets the FP8 clipping" ON)
 
-# Note: MIOPEN_FP8_IEEE_EXPONENT_BIAS can be defaulted by
-# cmake/MiopenCkCompileDefs.cmake to stay aligned with composable_kernel FP8 mode.
-
 configure_file("${PROJECT_SOURCE_DIR}/include/miopen/config.h.in" "${PROJECT_BINARY_DIR}/include/miopen/config.h")
 
 # configure a header file to pass the CMake version settings to the source, and package the header files in the output archive

@@ -30,6 +30,7 @@
 #include <miopen/logger.hpp>
 #include <miopen/tensor.hpp>
 #include <miopen/datatype.hpp>
+#include <miopen/target_properties.hpp>
 
 MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_CHECK_NUMERICS)
 
@@ -108,7 +109,8 @@ bool checkNumericsImpl(
     const std::vector<size_t> vgd = {numBlocks, size_t{1}, size_t{1}};
     const std::string build_options =
         " -DMIOPEN_FP8_CLIPPING=" + std::to_string(MIOPEN_FP8_CLIPPING) +
-        " -DMIOPEN_FP8_IEEE_EXPONENT_BIAS=" + std::to_string(MIOPEN_FP8_IEEE_EXPONENT_BIAS);
+        " -DMIOPEN_FP8_IEEE_EXPONENT_BIAS=" +
+        std::to_string(UsesIeeeExponentBias(handle.GetDeviceName()) ? 1 : 0);
     handle.AddKernel("MIOpenCheckNumerics",
                      "MIOpenCheckNumerics",
                      program_name,

@@ -354,7 +354,8 @@ void ProblemDescription::SetupFloats(ExecutionContext& ctx) const
 {
     if(IsFp32() || IsFp16() || IsBfp16() || IsInt8() || IsFp8() || IsBfp8())
     {
-        ctx.general_compile_options += GetDataTypeKernelParams(GetInDataType());
+        ctx.general_compile_options +=
+            GetDataTypeKernelParams(GetInDataType(), ctx.GetStream().GetDeviceName());
         return;
     }
 

@@ -75,10 +75,10 @@ std::string RefKernelName(const std::string& direction,
 }
 
 // Build compile options for the GPU reference kernel.
-std::string RefCompileOptions(miopenDataType_t data_type)
+std::string RefCompileOptions(const Handle& handle, miopenDataType_t data_type)
 {
     std::ostringstream ss;
-    ss << GetDataTypeKBP(data_type).GenerateFor(kbp::HIP{});
+    ss << GetDataTypeKBP(data_type, handle.GetDeviceName()).GenerateFor(kbp::HIP{});
     ss << " -DNAIVE_CONV_BLOCK_SIZE=" << REF_BLOCK_SIZE;
     return ss.str();
 }
@@ -265,7 +265,7 @@ void GpuConvReference::RunFwd(const Handle& handle,
 
         std::string kernel_name = RefKernelName("fwd", layout_str, data_type, yDesc.GetType());
         std::string kernel_file = "naive_conv_fwd.cpp";
-        std::string comp_opts   = RefCompileOptions(data_type);
+        std::string comp_opts   = RefCompileOptions(handle, data_type);
 
         auto in_strides =
             MakeStrideArray<5>(SplitStrideCtoGC(group, xDesc.GetStrides(), G_stride_idx));
@@ -346,7 +346,7 @@ void GpuConvReference::RunFwd(const Handle& handle,
 
         std::string kernel_name = RefKernelName("fwd", layout_str, data_type, yDesc.GetType());
         std::string kernel_file = "naive_conv_fwd.cpp";
-        std::string comp_opts   = RefCompileOptions(data_type);
+        std::string comp_opts   = RefCompileOptions(handle, data_type);
 
         auto in_strides =
             MakeStrideArray<6>(SplitStrideCtoGC(group, xDesc.GetStrides(), G_stride_idx));
@@ -467,7 +467,7 @@ void GpuConvReference::RunBwd(const Handle& handle,
 
         std::string kernel_name = RefKernelName("bwd", layout_str, data_type, dxDesc.GetType());
         std::string kernel_file = "naive_conv_bwd.cpp";
-        std::string comp_opts   = RefCompileOptions(data_type);
+        std::string comp_opts   = RefCompileOptions(handle, data_type);
 
         // BWD strides: dy is "out", dx is "in" from the kernel's perspective
         auto out_strides =
@@ -539,7 +539,7 @@ void GpuConvReference::RunBwd(const Handle& handle,
 
         std::string kernel_name = RefKernelName("bwd", layout_str, data_type, dxDesc.GetType());
         std::string kernel_file = "naive_conv_bwd.cpp";
-        std::string comp_opts   = RefCompileOptions(data_type);
+        std::string comp_opts   = RefCompileOptions(handle, data_type);
 
         auto out_strides =
             MakeStrideArray<6>(SplitStrideCtoGC(group, dyDesc.GetStrides(), G_stride_idx));
@@ -637,7 +637,7 @@ void GpuConvReference::RunWrw(const Handle& handle,
 
         std::string kernel_name = RefKernelName("wrw", layout_str, data_type, dwDesc.GetType());
         std::string kernel_file = "naive_conv_wrw.cpp";
-        std::string comp_opts   = RefCompileOptions(data_type);
+        std::string comp_opts   = RefCompileOptions(handle, data_type);
 
         auto in_strides =
             MakeStrideArray<5>(SplitStrideCtoGC(group, xDesc.GetStrides(), G_stride_idx));
@@ -697,7 +697,7 @@ void GpuConvReference::RunWrw(const Handle& handle,
 
         std::string kernel_name = RefKernelName("wrw", layout_str, data_type, dwDesc.GetType());
         std::string kernel_file = "naive_conv_wrw.cpp";
-        std::string comp_opts   = RefCompileOptions(data_type);
+        std::string comp_opts   = RefCompileOptions(handle, data_type);
 
         auto in_strides =
             MakeStrideArray<6>(SplitStrideCtoGC(group, xDesc.GetStrides(), G_stride_idx));

@@ -177,7 +177,7 @@ float Im2d2ColGPU(const Handle& handle,
         add_params(" -DNUM_IM_BLKS_EQ_1=" + std::to_string(static_cast<int>(num_blks == 1)));
         add_params(" -DUSE_IM_OFF_GUARD=1"); // always one
 
-        params += GetDataTypeKernelParams(type);
+        params += GetDataTypeKernelParams(type, handle.GetDeviceName());
 
         add_params(" -DLAYOUT_NHWC=" + std::to_string(static_cast<int>(layoutNHWC)));
         const int group_size_x = 256;
@@ -556,7 +556,7 @@ float Im3d2ColGPU(const Handle& handle,
     }
     else
     {
-        std::string params = GetDataTypeKernelParams(type);
+        std::string params = GetDataTypeKernelParams(type, handle.GetDeviceName());
         auto add_params    = [&](std::string param) {
             params += param;
             network_config += param;
@@ -671,7 +671,7 @@ float Col2Im2dGPU(const Handle& handle,
     }
     else
     {
-        std::string params = GetDataTypeKernelParams(type);
+        std::string params = GetDataTypeKernelParams(type, handle.GetDeviceName());
 
         size_t global_threads = static_cast<size_t>(in_c) * in_h * in_w;
         size_t local_threads  = std::min(WG_SIZE, global_threads);
@@ -806,7 +806,7 @@ float Col2Im3dGPU(const Handle& handle,
 
         const bool use_64_bit_index = index_size > 0xffffffffULL;
 
-        std::string params = GetDataTypeKernelParams(type);
+        std::string params = GetDataTypeKernelParams(type, handle.GetDeviceName());
 
         params += use_64_bit_index ? " -DMIOPEN_USE_64BIT_INDEX=1" : " -DMIOPEN_USE_64BIT_INDEX=0";
         params += " -DLAYOUT_NHWC=" + std::to_string(static_cast<int>(layoutNHWC));
@@ -939,7 +939,7 @@ float Col2Im3dGPUBatched(const Handle& handle,
 
         const bool use_64_bit_index = index_size > 0xffffffffULL;
 
-        std::string params = GetDataTypeKernelParams(type);
+        std::string params = GetDataTypeKernelParams(type, handle.GetDeviceName());
         params += use_64_bit_index ? " -DMIOPEN_USE_64BIT_INDEX=1" : " -DMIOPEN_USE_64BIT_INDEX=0";
         // Force a distinct code-object cache key for the batched entry point.
         params += " -DMIOPEN_COL2IM3D_BATCHED_VARIANT=1";
@@ -1157,7 +1157,7 @@ float transpose_NCHW2CNHW(const Handle& handle,
 
     std::string kernel_name = "transpose_NCHW2CNHW";
 
-    std::string params = GetDataTypeKernelParams(type);
+    std::string params = GetDataTypeKernelParams(type, handle.GetDeviceName());
 
     if(h_stride == 1 && w_stride == 1 && type == miopenFloat)
     {
@@ -1287,7 +1287,7 @@ float transpose_CNHW2NCHW(const Handle& handle,
 
     std::string kernel_name = "transpose_CNHW2NCHW";
 
-    std::string params = GetDataTypeKernelParams(type);
+    std::string params = GetDataTypeKernelParams(type, handle.GetDeviceName());
 
     if(h_stride == 1 && w_stride == 1 && type == miopenFloat)
     {
@@ -1549,7 +1549,7 @@ float transpose_packed_MN2NM(const Handle& handle,
 
     auto&& kernels = handle.GetKernels(kernel_name, network_config);
 
-    std::string params = GetDataTypeKernelParams(type);
+    std::string params = GetDataTypeKernelParams(type, handle.GetDeviceName());
 
     if(type != miopenInt8)
     {

@@ -28,6 +28,7 @@
 
 #include <miopen/errors.hpp>
 #include <miopen/kernel_build_params.hpp>
+#include <miopen/target_properties.hpp>
 #include <miopen/visit_float.hpp>
 
 #include <sstream>
@@ -131,7 +132,7 @@ inline std::size_t get_index_max(miopenIndexType_t index_type)
     MIOPEN_THROW("not belong to any case");
 }
 
-inline KernelBuildParameters GetDataTypeKBP(miopenDataType_t type)
+inline KernelBuildParameters GetDataTypeKBP(miopenDataType_t type, const std::string& device_name)
 {
     // values for MIOPEN_USE_ macros
     int use_fp16               = 0;
@@ -170,7 +171,7 @@ inline KernelBuildParameters GetDataTypeKBP(miopenDataType_t type)
         {"MIOPEN_USE_BFP16", use_bfp16},
         {"MIOPEN_USE_INT32", use_int32},
         {"MIOPEN_USE_RNE_BFLOAT16", use_rne_bfloat16},
-        {"MIOPEN_FP8_IEEE_EXPONENT_BIAS", MIOPEN_FP8_IEEE_EXPONENT_BIAS},
+        {"MIOPEN_FP8_IEEE_EXPONENT_BIAS", UsesIeeeExponentBias(device_name) ? 1 : 0},
         {"MIOPEN_FP8_CLIPPING", MIOPEN_FP8_CLIPPING},
     };
     if(use_fp64 != 0)
@@ -182,9 +183,9 @@ inline KernelBuildParameters GetDataTypeKBP(miopenDataType_t type)
     return kbp;
 }
 
-inline std::string GetDataTypeKernelParams(miopenDataType_t type)
+inline std::string GetDataTypeKernelParams(miopenDataType_t type, const std::string& device_name)
 {
-    return " " + GetDataTypeKBP(type).GenerateFor(kbp::HIP{});
+    return " " + GetDataTypeKBP(type, device_name).GenerateFor(kbp::HIP{});
 }
 
 } // namespace miopen

@@ -42,6 +42,12 @@ static inline bool IsTF32Supported(const std::string& device_name)
     return device_name == "gfx942" || StartsWith(device_name, "gfx95");
 }
 
+// Keep in sync with MIOPEN_CK_USE_OCP_FP8_TARGETS in cmake/MiopenCkCompileDefs.cmake.
+static inline bool UsesIeeeExponentBias(const std::string& device_name)
+{
+    return device_name == "gfx950" || StartsWith(device_name, "gfx12");
+}
+
 struct Handle;
 
 class TargetProperties

@@ -166,7 +166,8 @@ struct UniversalTransposeSolver : TransposePseudoSolver
         auto create_kernel = [&](const std::string& index_type) {
             const auto tensor_space = problem.input.GetElementSize();
             const auto cus          = ctx.GetStream().GetMaxComputeUnits();
-            auto build_params       = GetDataTypeKBP(problem.input.GetType());
+            auto build_params =
+                GetDataTypeKBP(problem.input.GetType(), ctx.GetStream().GetDeviceName());
 
             constexpr std::size_t max_block_size = 256;
             const auto local_size                = max_block_size;
@@ -268,9 +269,10 @@ struct TiledTransposeSolver : TransposePseudoSolver
         auto sln = ConvSolution{};
 
         auto create_kernel = [&](const std::string& index_type) {
-            const auto& lens  = problem.input.GetLengths();
-            const auto cus    = ctx.GetStream().GetMaxComputeUnits();
-            auto build_params = GetDataTypeKBP(problem.input.GetType());
+            const auto& lens = problem.input.GetLengths();
+            const auto cus   = ctx.GetStream().GetMaxComputeUnits();
+            auto build_params =
+                GetDataTypeKBP(problem.input.GetType(), ctx.GetStream().GetDeviceName());
 
             const auto n = lens.size() >= 1 ? lens[0] : 1;
             const auto c = lens.size() >= 2 ? lens[1] : 1;
@@ -387,7 +389,8 @@ struct VectorizedTransposeSolver : TransposePseudoSolver
         auto create_kernel = [&](const std::string& index_type) {
             const auto tensor_space = problem.input.GetElementSize();
             const auto cus          = ctx.GetStream().GetMaxComputeUnits();
-            auto build_params       = GetDataTypeKBP(problem.input.GetType());
+            auto build_params =
+                GetDataTypeKBP(problem.input.GetType(), ctx.GetStream().GetDeviceName());
 
             const auto num_blocks =
                 std::max(std::size_t{1}, (tensor_space + BlockSize - 1) / BlockSize);