From d95f83088c5a514487d09400c1782450f15625ef Mon Sep 17 00:00:00 2001 From: James Newling Date: Tue, 27 Oct 2026 14:44:06 -0700 Subject: [PATCH 1/4] Fix hipBLASLt GEMM tests to use explicit FNUZ host types --- projects/miopen/test/gtest/hipblaslt_gemm.cpp | 58 ++++++++++++++----- 1 file changed, 45 insertions(+), 13 deletions(-) diff --git a/projects/miopen/test/gtest/hipblaslt_gemm.cpp b/projects/miopen/test/gtest/hipblaslt_gemm.cpp index 9239748e5305..d29e69d049a0 100644 --- a/projects/miopen/test/gtest/hipblaslt_gemm.cpp +++ b/projects/miopen/test/gtest/hipblaslt_gemm.cpp @@ -42,11 +42,13 @@ #include #include -#include +// MIOpen test helpers already define half as half_float::half. +#ifndef HIP_NO_HALF +#define HIP_NO_HALF +#endif +#include -using float16 = half_float::half; -using float8_fnuz = miopen_f8::hip_f8; -using bfloat8_fnuz = miopen_f8::hip_f8; +using float16 = half_float::half; using namespace miopen; @@ -55,6 +57,10 @@ using namespace miopen; namespace hipblaslt_gemm { +// The descriptors below request FNUZ regardless of MIOPEN_FP8_IEEE_EXPONENT_BIAS. +using Float8Fnuz = __hip_fp8_e4m3_fnuz; +using BFloat8Fnuz = __hip_fp8_e5m2_fnuz; + struct TestCase { bool isColMajor; @@ -189,7 +195,8 @@ static void RunGemmDescriptors(const TestCase& testCase, miopenDataType_t dataTy for(auto& index : workspaceB_host) { - index = prng::gen_A_to_B(static_cast(-0.5), static_cast(0.5f)); + // Quantize before shifting, preserving the input distribution for low precision types. + index = static_cast(static_cast(prng::gen_canonical()) - 0.5f); } workspaceA_device.Write(workspaceA_host); @@ -228,7 +235,11 @@ static void RunGemmDescriptors(const TestCase& testCase, miopenDataType_t dataTy miopen::gemm_cpu_util::CallGemm( desc, workspaceA_host.data(), 0, workspaceB_host.data(), 0, workspaceC_host.data(), 0); - auto error = miopen::rms_range(workspaceC_host, workspaceC_device.Read>()); + const auto result = workspaceC_device.Read>(); + const auto referenceFloat = + std::vector(workspaceC_host.begin(), workspaceC_host.end()); + const auto resultFloat = std::vector(result.begin(), result.end()); + auto error = miopen::rms_range(referenceFloat, resultFloat); EXPECT_TRUE(std::isfinite(error)); const double tolerance = ((sizeof(T) == 4) ? 1e-6 : 7e-2); @@ -288,6 +299,28 @@ static void CheckExceptionsWithSkip([[maybe_unused]] miopenDataType_t dataType) } // namespace hipblaslt_gemm using namespace hipblaslt_gemm; +TEST(CPU_HipBLASLtGEMMTest_FP8, FNUZEncoding) +{ + EXPECT_EQ(Float8Fnuz(1.0f).__x, 0x40); + EXPECT_EQ(Float8Fnuz(-0.0f).__x, 0x00); + Float8Fnuz value; + value.__x = 0x40; + EXPECT_EQ(static_cast(value), 1.0f); + value.__x = 0x80; + EXPECT_TRUE(std::isnan(static_cast(value))); +} + +TEST(CPU_HipBLASLtGEMMTest_BFP8, FNUZEncoding) +{ + EXPECT_EQ(BFloat8Fnuz(1.0f).__x, 0x40); + EXPECT_EQ(BFloat8Fnuz(-0.0f).__x, 0x00); + BFloat8Fnuz value; + value.__x = 0x40; + EXPECT_EQ(static_cast(value), 1.0f); + value.__x = 0x80; + EXPECT_TRUE(std::isnan(static_cast(value))); +} + TEST_F(GPU_HipBLASLtGEMMTest_FP32, CheckHipBLASLtGEMMException) { using e_mask = enabled; @@ -340,7 +373,7 @@ TEST_F(GPU_HipBLASLtGEMMTest_FP8, CheckHipBLASLtGEMMException) Gpu::gfx90A, Gpu::gfx110X, Gpu::gfx115X>; - CheckExceptionsWithSkip(miopenDataType_t::miopenFloat8_fnuz); + CheckExceptionsWithSkip(miopenDataType_t::miopenFloat8_fnuz); }; TEST_P(GPU_HipBLASLtGEMMTest_FP8, RunHipBLASLtGEMM) { @@ -352,8 +385,7 @@ TEST_P(GPU_HipBLASLtGEMMTest_FP8, RunHipBLASLtGEMM) Gpu::gfx90A, Gpu::gfx110X, Gpu::gfx115X>; - RunGemmDescriptors(GetParam(), - miopenDataType_t::miopenFloat8_fnuz); + RunGemmDescriptors(GetParam(), miopenDataType_t::miopenFloat8_fnuz); }; INSTANTIATE_TEST_SUITE_P(Full, GPU_HipBLASLtGEMMTest_FP8, testing::ValuesIn(GetTestCases())); @@ -368,9 +400,9 @@ TEST_F(GPU_HipBLASLtGEMMTest_BFP8, CheckHipBLASLtGEMMException) Gpu::gfx90A, Gpu::gfx110X, Gpu::gfx115X>; - CheckExceptionsWithSkip(miopenDataType_t::miopenBFloat8_fnuz); + CheckExceptionsWithSkip(miopenDataType_t::miopenBFloat8_fnuz); #else - CheckExceptions(miopenDataType_t::miopenInt64); + CheckExceptions(miopenDataType_t::miopenInt64); #endif }; TEST_P(GPU_HipBLASLtGEMMTest_BFP8, RunHipBLASLtGEMM) @@ -384,8 +416,8 @@ TEST_P(GPU_HipBLASLtGEMMTest_BFP8, RunHipBLASLtGEMM) Gpu::gfx90A, Gpu::gfx110X, Gpu::gfx115X>; - RunGemmDescriptors(GetParam(), - miopenDataType_t::miopenBFloat8_fnuz); + RunGemmDescriptors(GetParam(), + miopenDataType_t::miopenBFloat8_fnuz); #else GTEST_SKIP(); #endif From 2415d1774b4dbbe25afd43f5d9948f5d5161cc50 Mon Sep 17 00:00:00 2001 From: Yassin Soliman Date: Wed, 28 Oct 2026 15:39:54 -0600 Subject: [PATCH 2/4] fix(miopen): pick the FP8 exponent bias from the device at run time MIOpen chose MIOPEN_FP8_IEEE_EXPONENT_BIAS once at configure time from GPU_TARGETS. A host build for several targets that included gfx950 or gfx12 therefore used the IEEE bias on gfx942 as well, which broke the FP8 and BF8 hipBLASLt GEMM tests there (issue #14375). - Add miopen::UsesIeeeExponentBias(device_name): true for gfx950 and gfx12, matching MIOPEN_CK_USE_OCP_FP8_TARGETS. - GetDataTypeKBP and GetDataTypeKernelParams take the device name and emit the bias for that device. Update every caller. - check_numerics, the naive conv solver and the MHA solvers use the device helper instead of the config.h macro. - Remove the CMake seed, the MIOPEN_FP8_IEEE_EXPONENT_BIAS option and its config.h entry. - Host FP8 types in hip_float8.hpp read the bias from a run-time setting. MIOpenDriver and the FP8 solver test fixture set it from the device. --- projects/miopen/CHANGELOG.md | 6 ++ .../miopen/cmake/MiopenCkCompileDefs.cmake | 18 ------ projects/miopen/driver/driver.hpp | 4 ++ projects/miopen/include/miopen/config.h.in | 1 - projects/miopen/src/CMakeLists.txt | 4 -- projects/miopen/src/check_numerics.cpp | 4 +- .../miopen/src/conv/problem_description.cpp | 3 +- projects/miopen/src/gpu_conv_reference.cpp | 16 ++--- projects/miopen/src/hip/utilocl.cpp | 16 ++--- .../miopen/src/include/miopen/datatype.hpp | 9 +-- .../src/include/miopen/target_properties.hpp | 6 ++ .../miopen/utility/transposing_solver.hpp | 13 ++-- projects/miopen/src/kernels/hip_float8.hpp | 59 +++++++++++++------ projects/miopen/src/solver/adam/adam.cpp | 3 +- .../src/solver/adam/transformers_adam_w.cpp | 3 +- .../solver/conv/conv_direct_naive_conv.cpp | 4 +- .../src/solver/mha/mha_solver_backward.cpp | 58 +++++++++--------- .../src/solver/mha/mha_solver_forward.cpp | 38 ++++++------ .../miopen/src/solver/pooling/backward2d.cpp | 4 +- .../miopen/src/solver/pooling/backwardNd.cpp | 4 +- .../miopen/src/solver/pooling/forward2d.cpp | 5 +- .../src/solver/pooling/forwardNaive.cpp | 4 +- .../miopen/src/solver/pooling/forwardNd.cpp | 5 +- projects/miopen/src/tensor.cpp | 10 ++-- projects/miopen/test/gtest/solver_f8.hpp | 2 + 25 files changed, 166 insertions(+), 133 deletions(-) diff --git a/projects/miopen/CHANGELOG.md b/projects/miopen/CHANGELOG.md index cd0cbfe89c1f..25880e7d5cff 100644 --- a/projects/miopen/CHANGELOG.md +++ b/projects/miopen/CHANGELOG.md @@ -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 diff --git a/projects/miopen/cmake/MiopenCkCompileDefs.cmake b/projects/miopen/cmake/MiopenCkCompileDefs.cmake index 59d7e51ac191..63fca1f61177 100644 --- a/projects/miopen/cmake/MiopenCkCompileDefs.cmake +++ b/projects/miopen/cmake/MiopenCkCompileDefs.cmake @@ -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() diff --git a/projects/miopen/driver/driver.hpp b/projects/miopen/driver/driver.hpp index e16ce9b18f04..1735693cb759 100644 --- a/projects/miopen/driver/driver.hpp +++ b/projects/miopen/driver/driver.hpp @@ -551,11 +551,15 @@ template <> inline void Driver::InitDataType() { data_type = miopenFloat8_fnuz; + miopen_f8::set_hip_f8_ieee_exponent_bias( + miopen::UsesIeeeExponentBias(miopen::deref(GetHandle()).GetDeviceName())); } template <> inline void Driver::InitDataType() { data_type = miopenBFloat8_fnuz; + miopen_f8::set_hip_f8_ieee_exponent_bias( + miopen::UsesIeeeExponentBias(miopen::deref(GetHandle()).GetDeviceName())); } // "std::is_same{}" used to avoid "static_assert" compilation error, // which occurs when the condition does not depend in any way on the template parameters. diff --git a/projects/miopen/include/miopen/config.h.in b/projects/miopen/include/miopen/config.h.in index 658776cc46a9..2a60389d0927 100644 --- a/projects/miopen/include/miopen/config.h.in +++ b/projects/miopen/include/miopen/config.h.in @@ -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 diff --git a/projects/miopen/src/CMakeLists.txt b/projects/miopen/src/CMakeLists.txt index 1b7d95f11ea2..668185c3ff98 100644 --- a/projects/miopen/src/CMakeLists.txt +++ b/projects/miopen/src/CMakeLists.txt @@ -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 diff --git a/projects/miopen/src/check_numerics.cpp b/projects/miopen/src/check_numerics.cpp index b75f7ad4d057..dd586bb3b965 100644 --- a/projects/miopen/src/check_numerics.cpp +++ b/projects/miopen/src/check_numerics.cpp @@ -30,6 +30,7 @@ #include #include #include +#include MIOPEN_DECLARE_ENV_VAR_UINT64(MIOPEN_CHECK_NUMERICS) @@ -108,7 +109,8 @@ bool checkNumericsImpl( const std::vector 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, diff --git a/projects/miopen/src/conv/problem_description.cpp b/projects/miopen/src/conv/problem_description.cpp index a06cc8ba6632..d00ee3332820 100644 --- a/projects/miopen/src/conv/problem_description.cpp +++ b/projects/miopen/src/conv/problem_description.cpp @@ -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; } diff --git a/projects/miopen/src/gpu_conv_reference.cpp b/projects/miopen/src/gpu_conv_reference.cpp index 83db66d1a7e4..07f759391f87 100644 --- a/projects/miopen/src/gpu_conv_reference.cpp +++ b/projects/miopen/src/gpu_conv_reference.cpp @@ -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)); diff --git a/projects/miopen/src/hip/utilocl.cpp b/projects/miopen/src/hip/utilocl.cpp index 77e3b564a88c..923e4dfca3a6 100644 --- a/projects/miopen/src/hip/utilocl.cpp +++ b/projects/miopen/src/hip/utilocl.cpp @@ -177,7 +177,7 @@ float Im2d2ColGPU(const Handle& handle, add_params(" -DNUM_IM_BLKS_EQ_1=" + std::to_string(static_cast(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(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(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(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) { diff --git a/projects/miopen/src/include/miopen/datatype.hpp b/projects/miopen/src/include/miopen/datatype.hpp index 74a7a1d95074..76208e6de941 100644 --- a/projects/miopen/src/include/miopen/datatype.hpp +++ b/projects/miopen/src/include/miopen/datatype.hpp @@ -28,6 +28,7 @@ #include #include +#include #include #include @@ -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 diff --git a/projects/miopen/src/include/miopen/target_properties.hpp b/projects/miopen/src/include/miopen/target_properties.hpp index e79fdf9bfb94..5e8af5b42fa5 100644 --- a/projects/miopen/src/include/miopen/target_properties.hpp +++ b/projects/miopen/src/include/miopen/target_properties.hpp @@ -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 diff --git a/projects/miopen/src/include/miopen/utility/transposing_solver.hpp b/projects/miopen/src/include/miopen/utility/transposing_solver.hpp index 2fe4a4dd1827..ae162dad6dd4 100644 --- a/projects/miopen/src/include/miopen/utility/transposing_solver.hpp +++ b/projects/miopen/src/include/miopen/utility/transposing_solver.hpp @@ -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); diff --git a/projects/miopen/src/kernels/hip_float8.hpp b/projects/miopen/src/kernels/hip_float8.hpp index 412ae7846ce5..6fdafe8865c2 100644 --- a/projects/miopen/src/kernels/hip_float8.hpp +++ b/projects/miopen/src/kernels/hip_float8.hpp @@ -46,9 +46,9 @@ #define USE_SIMPLER_HIP_F8x8 0 -// FP8 configuration macros. In host builds these come from config.h -// (included transitively). In kernel builds the solver passes them -// via -D compiler flags. Fallback defaults match CMake defaults. +// FP8 configuration macros. Kernel builds get them as -D compiler flags from the solver, which +// picks the exponent bias for the target device. Host code starts from these defaults and can +// switch the exponent bias at run time with set_hip_f8_ieee_exponent_bias(). #ifndef MIOPEN_FP8_CLIPPING #define MIOPEN_FP8_CLIPPING 1 #endif @@ -82,15 +82,40 @@ enum class hip_f8_rounding_mode stochastic }; -inline MIOPEN_HIP_HOST_DEVICE bool get_hip_f8_bias_mode() +#if(defined(__HIP_DEVICE_COMPILE__) && __HIP_DEVICE_COMPILE__) || defined(__HIPCC_RTC__) +#define MIOPEN_FP8_HOST_RUNTIME_FORMAT 0 +#else +#define MIOPEN_FP8_HOST_RUNTIME_FORMAT 1 +#endif + +#ifndef __HIPCC_RTC__ +// One host binary can drive devices that use different FP8 formats, so the host picks the +// exponent bias at run time. Set it from miopen::UsesIeeeExponentBias(device name) before +// converting values for a device. Declared in device passes too, so host code in HIP sources +// compiles there, but device code reads the -D value. +inline bool& host_ieee_exponent_bias() +{ + static bool value = MIOPEN_FP8_IEEE_EXPONENT_BIAS != 0; + return value; +} + +inline void set_hip_f8_ieee_exponent_bias(bool ieee) { host_ieee_exponent_bias() = ieee; } +#endif + +inline MIOPEN_HIP_HOST_DEVICE bool get_hip_f8_ieee_exponent_bias() { -#if MIOPEN_FP8_IEEE_EXPONENT_BIAS - return false; +#if MIOPEN_FP8_HOST_RUNTIME_FORMAT + return host_ieee_exponent_bias(); #else - return true; + return MIOPEN_FP8_IEEE_EXPONENT_BIAS != 0; #endif } +inline MIOPEN_HIP_HOST_DEVICE bool get_hip_f8_bias_mode() +{ + return !get_hip_f8_ieee_exponent_bias(); +} + template class numeric_limits; @@ -482,12 +507,10 @@ inline MIOPEN_HIP_HOST_DEVICE bool operator>(const miopen_f8::hip_f8& lhs, template inline MIOPEN_HIP_HOST_DEVICE miopen_f8::hip_f8 fabs(miopen_f8::hip_f8 v) { -#if !MIOPEN_FP8_IEEE_EXPONENT_BIAS // Preserve "universal" nan. // Otherwise it will be converted to valid 0. - if(v.data == 0x80) + if(!miopen_f8::get_hip_f8_ieee_exponent_bias() && v.data == 0x80) return v; -#endif v.data = v.data & 0x7f; return v; } @@ -512,25 +535,25 @@ class numeric_limits> static MIOPEN_HIP_HOST_DEVICE miopen_f8::hip_f8 epsilon() { return miopen_f8::Generate>( - MIOPEN_FP8_IEEE_EXPONENT_BIAS ? 0x20 : 0x28); // 0.125 + miopen_f8::get_hip_f8_ieee_exponent_bias() ? 0x20 : 0x28); // 0.125 } static MIOPEN_HIP_HOST_DEVICE miopen_f8::hip_f8 quiet_NaN() { return miopen_f8::Generate>( - MIOPEN_FP8_IEEE_EXPONENT_BIAS ? 0x7c : 0x80); + miopen_f8::get_hip_f8_ieee_exponent_bias() ? 0x7c : 0x80); } static MIOPEN_HIP_HOST_DEVICE miopen_f8::hip_f8 signaling_NaN() { return miopen_f8::Generate>( - MIOPEN_FP8_IEEE_EXPONENT_BIAS ? 0x79 : 0x80); + miopen_f8::get_hip_f8_ieee_exponent_bias() ? 0x79 : 0x80); } static MIOPEN_HIP_HOST_DEVICE miopen_f8::hip_f8 max() { return miopen_f8::Generate>( - MIOPEN_FP8_IEEE_EXPONENT_BIAS ? 0x77 : 0x7f); // 240 + miopen_f8::get_hip_f8_ieee_exponent_bias() ? 0x77 : 0x7f); // 240 } static MIOPEN_HIP_HOST_DEVICE miopen_f8::hip_f8 min() @@ -553,25 +576,25 @@ class numeric_limits> static MIOPEN_HIP_HOST_DEVICE miopen_f8::hip_f8 epsilon() { return miopen_f8::Generate>( - MIOPEN_FP8_IEEE_EXPONENT_BIAS ? 0x34 : 0x38); // 0.25 + miopen_f8::get_hip_f8_ieee_exponent_bias() ? 0x34 : 0x38); // 0.25 } static MIOPEN_HIP_HOST_DEVICE miopen_f8::hip_f8 quiet_NaN() { return miopen_f8::Generate>( - MIOPEN_FP8_IEEE_EXPONENT_BIAS ? 0x7e : 0x80); + miopen_f8::get_hip_f8_ieee_exponent_bias() ? 0x7e : 0x80); } static MIOPEN_HIP_HOST_DEVICE miopen_f8::hip_f8 signaling_NaN() { return miopen_f8::Generate>( - MIOPEN_FP8_IEEE_EXPONENT_BIAS ? 0x7d : 0x80); + miopen_f8::get_hip_f8_ieee_exponent_bias() ? 0x7d : 0x80); } static MIOPEN_HIP_HOST_DEVICE miopen_f8::hip_f8 max() { return miopen_f8::Generate>( - MIOPEN_FP8_IEEE_EXPONENT_BIAS ? 0x7b : 0x7f); // 57344 + miopen_f8::get_hip_f8_ieee_exponent_bias() ? 0x7b : 0x7f); // 57344 } static MIOPEN_HIP_HOST_DEVICE miopen_f8::hip_f8 min() diff --git a/projects/miopen/src/solver/adam/adam.cpp b/projects/miopen/src/solver/adam/adam.cpp index 0b39bf5bf358..3d5fd8126685 100644 --- a/projects/miopen/src/solver/adam/adam.cpp +++ b/projects/miopen/src/solver/adam/adam.cpp @@ -64,7 +64,8 @@ ConvSolution Adam::GetSolution(const ExecutionContext& context, {"GTYPE", grad_dtype}, {"CTYPE", ptype_size > 4 ? "double" : "float"}, } - << GetDataTypeKBP(problem.GetParamDesc().GetType()); + << GetDataTypeKBP(problem.GetParamDesc().GetType(), + context.GetStream().GetDeviceName()); constexpr size_t local_size = 256; auto& handle = context.GetStream(); diff --git a/projects/miopen/src/solver/adam/transformers_adam_w.cpp b/projects/miopen/src/solver/adam/transformers_adam_w.cpp index 8eafdf7a9d81..c9f0033c4938 100644 --- a/projects/miopen/src/solver/adam/transformers_adam_w.cpp +++ b/projects/miopen/src/solver/adam/transformers_adam_w.cpp @@ -64,7 +64,8 @@ ConvSolution TransformersAdamW::GetSolution(const ExecutionContext& context, {"GTYPE", grad_dtype}, {"CTYPE", ptype_size > 4 ? "double" : "float"}, } - << GetDataTypeKBP(problem.GetParamDesc().GetType()); + << GetDataTypeKBP(problem.GetParamDesc().GetType(), + context.GetStream().GetDeviceName()); constexpr size_t local_size = 256; auto& handle = context.GetStream(); diff --git a/projects/miopen/src/solver/conv/conv_direct_naive_conv.cpp b/projects/miopen/src/solver/conv/conv_direct_naive_conv.cpp index 8025014dcfb1..3328c8f8808b 100644 --- a/projects/miopen/src/solver/conv/conv_direct_naive_conv.cpp +++ b/projects/miopen/src/solver/conv/conv_direct_naive_conv.cpp @@ -33,6 +33,7 @@ #include #include #include +#include #include MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_DEBUG_CONV_DIRECT_NAIVE_USE_PACKED_KERNELS); @@ -285,7 +286,8 @@ std::string ConvDirectNaiveConvCompileOption(const ExecutionContext& ctx, if(out_cast_type) ss << " -DOUTPUT_CAST_TYPE=" << miopen::GetDataType(*out_cast_type); ss << " -DMIOPEN_FP8_CLIPPING=" << MIOPEN_FP8_CLIPPING; - ss << " -DMIOPEN_FP8_IEEE_EXPONENT_BIAS=" << MIOPEN_FP8_IEEE_EXPONENT_BIAS; + ss << " -DMIOPEN_FP8_IEEE_EXPONENT_BIAS=" + << (UsesIeeeExponentBias(ctx.GetStream().GetDeviceName()) ? 1 : 0); // Let the kernel choose its accumulator (double for naive kernels ) } diff --git a/projects/miopen/src/solver/mha/mha_solver_backward.cpp b/projects/miopen/src/solver/mha/mha_solver_backward.cpp index 1551b28e1246..668fda650b1a 100644 --- a/projects/miopen/src/solver/mha/mha_solver_backward.cpp +++ b/projects/miopen/src/solver/mha/mha_solver_backward.cpp @@ -94,35 +94,35 @@ bool MhaBackward::IsApplicable( auto [N, H, S, D] = miopen::tien<4>(descsBwd.kDesc.GetLengths()); - return !env::disabled(MIOPEN_DEBUG_ATTN_NAIVE_BWD) // - && S <= std::numeric_limits::max() // - && D <= std::numeric_limits::max() // - && descsBwd.kDesc.IsPacked() // - && descsBwd.qDesc.IsPacked() // - && descsBwd.vDesc.IsPacked() // - && descsBwd.oDesc.IsPacked() // - && descsBwd.doDesc.IsPacked() // - && descsBwd.mDesc.IsPacked() // - && descsBwd.zInvDesc.IsPacked() // - && descsBwd.dkDesc.IsPacked() // - && descsBwd.dqDesc.IsPacked() // - && descsBwd.dvDesc.IsPacked() // - && descsBwd.mDesc.GetType() == miopenFloat // - && descsBwd.zInvDesc.GetType() == miopenFloat // - && descsBwd.kDesc.GetType() == descsBwd.qDesc.GetType() // - && descsBwd.kDesc.GetType() == descsBwd.vDesc.GetType() // - && descsBwd.kDesc.GetType() == descsBwd.oDesc.GetType() // - && descsBwd.kDesc.GetType() == descsBwd.dqDesc.GetType() // - && descsBwd.kDesc.GetType() == descsBwd.dkDesc.GetType() // - && descsBwd.kDesc.GetType() == descsBwd.dvDesc.GetType() // - && ((descsBwd.kDesc.GetType() == miopenFloat) // - || (USE_ROCBLAS_EX3 // - && (MIOPEN_FP8_IEEE_EXPONENT_BIAS == 0) // - && (descsBwd.kDesc.GetType() == miopenFloat8_fnuz))) // - && ((descsBwd.doDesc.GetType() == miopenFloat) // - || (USE_ROCBLAS_EX3 // - && (MIOPEN_FP8_IEEE_EXPONENT_BIAS == 0) // - && (descsBwd.doDesc.GetType() == miopenBFloat8_fnuz))); // + return !env::disabled(MIOPEN_DEBUG_ATTN_NAIVE_BWD) // + && S <= std::numeric_limits::max() // + && D <= std::numeric_limits::max() // + && descsBwd.kDesc.IsPacked() // + && descsBwd.qDesc.IsPacked() // + && descsBwd.vDesc.IsPacked() // + && descsBwd.oDesc.IsPacked() // + && descsBwd.doDesc.IsPacked() // + && descsBwd.mDesc.IsPacked() // + && descsBwd.zInvDesc.IsPacked() // + && descsBwd.dkDesc.IsPacked() // + && descsBwd.dqDesc.IsPacked() // + && descsBwd.dvDesc.IsPacked() // + && descsBwd.mDesc.GetType() == miopenFloat // + && descsBwd.zInvDesc.GetType() == miopenFloat // + && descsBwd.kDesc.GetType() == descsBwd.qDesc.GetType() // + && descsBwd.kDesc.GetType() == descsBwd.vDesc.GetType() // + && descsBwd.kDesc.GetType() == descsBwd.oDesc.GetType() // + && descsBwd.kDesc.GetType() == descsBwd.dqDesc.GetType() // + && descsBwd.kDesc.GetType() == descsBwd.dkDesc.GetType() // + && descsBwd.kDesc.GetType() == descsBwd.dvDesc.GetType() // + && ((descsBwd.kDesc.GetType() == miopenFloat) // + || (USE_ROCBLAS_EX3 // + && ((descsBwd.kDesc.GetType() == miopenFloat8_fnuz) // + && !UsesIeeeExponentBias(context.GetStream().GetDeviceName())))) // + && ((descsBwd.doDesc.GetType() == miopenFloat) // + || (USE_ROCBLAS_EX3 // + && ((descsBwd.doDesc.GetType() == miopenBFloat8_fnuz) // + && !UsesIeeeExponentBias(context.GetStream().GetDeviceName())))); // #else return false; #endif diff --git a/projects/miopen/src/solver/mha/mha_solver_forward.cpp b/projects/miopen/src/solver/mha/mha_solver_forward.cpp index ea62a1b244af..e7ddbfd21675 100644 --- a/projects/miopen/src/solver/mha/mha_solver_forward.cpp +++ b/projects/miopen/src/solver/mha/mha_solver_forward.cpp @@ -80,25 +80,25 @@ bool MhaForward::IsApplicable([[maybe_unused]] const ExecutionContext& context, auto [N, H, S, D] = miopen::tien<4>(descsFwd.kDesc.GetLengths()); - return !env::disabled(MIOPEN_DEBUG_ATTN_NAIVE_FWD) // - && S <= std::numeric_limits::max() // - && descsFwd.kDesc.IsPacked() // - && descsFwd.qDesc.IsPacked() // - && descsFwd.vDesc.IsPacked() // - && descsFwd.oDesc.IsPacked() // - && descsFwd.mDesc.IsPacked() // - && descsFwd.zInvDesc.IsPacked() // - && descsFwd.biasDesc.IsPacked() // - && descsFwd.mDesc.GetType() == miopenFloat // - && descsFwd.zInvDesc.GetType() == miopenFloat // - && descsFwd.biasDesc.GetType() == miopenFloat // - && descsFwd.kDesc.GetType() == descsFwd.qDesc.GetType() // - && descsFwd.kDesc.GetType() == descsFwd.vDesc.GetType() // - && descsFwd.kDesc.GetType() == descsFwd.oDesc.GetType() // - && ((descsFwd.kDesc.GetType() == miopenFloat) // - || (USE_ROCBLAS_EX3 // - && (MIOPEN_FP8_IEEE_EXPONENT_BIAS == 0) // - && (descsFwd.kDesc.GetType() == miopenFloat8_fnuz))); // + return !env::disabled(MIOPEN_DEBUG_ATTN_NAIVE_FWD) // + && S <= std::numeric_limits::max() // + && descsFwd.kDesc.IsPacked() // + && descsFwd.qDesc.IsPacked() // + && descsFwd.vDesc.IsPacked() // + && descsFwd.oDesc.IsPacked() // + && descsFwd.mDesc.IsPacked() // + && descsFwd.zInvDesc.IsPacked() // + && descsFwd.biasDesc.IsPacked() // + && descsFwd.mDesc.GetType() == miopenFloat // + && descsFwd.zInvDesc.GetType() == miopenFloat // + && descsFwd.biasDesc.GetType() == miopenFloat // + && descsFwd.kDesc.GetType() == descsFwd.qDesc.GetType() // + && descsFwd.kDesc.GetType() == descsFwd.vDesc.GetType() // + && descsFwd.kDesc.GetType() == descsFwd.oDesc.GetType() // + && ((descsFwd.kDesc.GetType() == miopenFloat) // + || (USE_ROCBLAS_EX3 // + && ((descsFwd.kDesc.GetType() == miopenFloat8_fnuz) // + && !UsesIeeeExponentBias(context.GetStream().GetDeviceName())))); // #else return false; #endif diff --git a/projects/miopen/src/solver/pooling/backward2d.cpp b/projects/miopen/src/solver/pooling/backward2d.cpp index 6c2ddbcd709d..14a5a6a09e69 100644 --- a/projects/miopen/src/solver/pooling/backward2d.cpp +++ b/projects/miopen/src/solver/pooling/backward2d.cpp @@ -141,7 +141,7 @@ bool PoolingBackward2d::IsApplicable(const ExecutionContext&, problem.GetYDesc().IsPossibleLayout4D5D("NCHW", strict); } -ConvSolution PoolingBackward2d::GetSolution(const ExecutionContext&, +ConvSolution PoolingBackward2d::GetSolution(const ExecutionContext& context, const miopen::pooling::ProblemDescription& problem, const PerformanceConfigPooling2dBackward& config) const { @@ -200,7 +200,7 @@ ConvSolution PoolingBackward2d::GetSolution(const ExecutionContext&, ? 1 : 0}, } - << GetDataTypeKBP(problem.GetXDesc().GetType()); + << GetDataTypeKBP(problem.GetXDesc().GetType(), context.GetStream().GetDeviceName()); kernel.comp_options = build_params.GenerateFor(kbp::HIP{}); diff --git a/projects/miopen/src/solver/pooling/backwardNd.cpp b/projects/miopen/src/solver/pooling/backwardNd.cpp index 7154126788df..8b0a2d1181e7 100644 --- a/projects/miopen/src/solver/pooling/backwardNd.cpp +++ b/projects/miopen/src/solver/pooling/backwardNd.cpp @@ -45,7 +45,7 @@ bool PoolingBackwardNd::IsApplicable(const ExecutionContext&, && problem.GetPooling().GetWorkspaceIndexMode() == miopenPoolingWorkspaceIndexMask); } -ConvSolution PoolingBackwardNd::GetSolution(const ExecutionContext&, +ConvSolution PoolingBackwardNd::GetSolution(const ExecutionContext& context, const miopen::pooling::ProblemDescription& problem, const PerformanceConfigPoolingNdBackward& config) const { @@ -146,7 +146,7 @@ ConvSolution PoolingBackwardNd::GetSolution(const ExecutionContext&, {"MLO_POOLING_INDEX_TYPE", get_pooling_index_type_name(problem.GetPooling().GetIndexType())}, } - << GetDataTypeKBP(problem.GetDYDesc().GetType()); + << GetDataTypeKBP(problem.GetDYDesc().GetType(), context.GetStream().GetDeviceName()); kernel.comp_options = build_params.GenerateFor(kbp::HIP{}); diff --git a/projects/miopen/src/solver/pooling/forward2d.cpp b/projects/miopen/src/solver/pooling/forward2d.cpp index bb1cd522e810..2e401db142da 100644 --- a/projects/miopen/src/solver/pooling/forward2d.cpp +++ b/projects/miopen/src/solver/pooling/forward2d.cpp @@ -138,7 +138,7 @@ bool PoolingForward2d::IsApplicable(const ExecutionContext& context, TargetProperties::GetMaxWaveScratchSize() / context.GetStream().GetWavefrontWidth(); } -ConvSolution PoolingForward2d::GetSolution(const ExecutionContext&, +ConvSolution PoolingForward2d::GetSolution(const ExecutionContext& context, const miopen::pooling::ProblemDescription& problem, const PerformanceConfigPooling2dForward& config) const { @@ -192,7 +192,8 @@ ConvSolution PoolingForward2d::GetSolution(const ExecutionContext&, }; } - build_params << GetDataTypeKBP(problem.GetXDesc().GetType()); + build_params << GetDataTypeKBP(problem.GetXDesc().GetType(), + context.GetStream().GetDeviceName()); kernel.comp_options = build_params.GenerateFor(kbp::HIP{}); diff --git a/projects/miopen/src/solver/pooling/forwardNaive.cpp b/projects/miopen/src/solver/pooling/forwardNaive.cpp index d745ae6e5c78..3c8868f5767c 100644 --- a/projects/miopen/src/solver/pooling/forwardNaive.cpp +++ b/projects/miopen/src/solver/pooling/forwardNaive.cpp @@ -71,7 +71,7 @@ bool PoolingForwardNaive::IsApplicable(const ExecutionContext&, } ConvSolution -PoolingForwardNaive::GetSolution(const ExecutionContext&, +PoolingForwardNaive::GetSolution(const ExecutionContext& context, const miopen::pooling::ProblemDescription& problem, const PerformanceConfigPoolingForwardNaive& config) const { @@ -214,7 +214,7 @@ PoolingForwardNaive::GetSolution(const ExecutionContext&, {"MLO_POOLING_INDEX_TYPE", get_pooling_index_type_name(index_type)}, {"MLO_POOLING_IS2D_KERNEL", static_cast(is2d_kernel)}, }; - build_params << GetDataTypeKBP(bot.GetType()); + build_params << GetDataTypeKBP(bot.GetType(), context.GetStream().GetDeviceName()); kernel.comp_options = build_params.GenerateFor(kbp::HIP{}); // [Informative] The total number of kernels required to cover the whole diff --git a/projects/miopen/src/solver/pooling/forwardNd.cpp b/projects/miopen/src/solver/pooling/forwardNd.cpp index cd17a956043b..14bb61b7600b 100644 --- a/projects/miopen/src/solver/pooling/forwardNd.cpp +++ b/projects/miopen/src/solver/pooling/forwardNd.cpp @@ -107,7 +107,7 @@ bool PoolingForwardNd::IsApplicable(const ExecutionContext& context, && problem.SaveIndex() == true); } -ConvSolution PoolingForwardNd::GetSolution(const ExecutionContext&, +ConvSolution PoolingForwardNd::GetSolution(const ExecutionContext& context, const miopen::pooling::ProblemDescription& problem, const PerformanceConfigPoolingNdForward& config) const { @@ -175,7 +175,8 @@ ConvSolution PoolingForwardNd::GetSolution(const ExecutionContext&, }; } - build_params << GetDataTypeKBP(problem.GetXDesc().GetType()); + build_params << GetDataTypeKBP(problem.GetXDesc().GetType(), + context.GetStream().GetDeviceName()); kernel.comp_options = build_params.GenerateFor(kbp::HIP{}); diff --git a/projects/miopen/src/tensor.cpp b/projects/miopen/src/tensor.cpp index 03933c828b2e..fa65cb84752a 100644 --- a/projects/miopen/src/tensor.cpp +++ b/projects/miopen/src/tensor.cpp @@ -1117,7 +1117,7 @@ void SetTensor(const Handle& handle, std::stringstream ss; // SUBTENSOR_OP_WITH_SCALAR set to 0 for set operation, and 1 for multiply operation ss << "-DSUBTENSOR_OP_WITH_SCALAR=0"; - ss << GetDataTypeKernelParams(dataType); + ss << GetDataTypeKernelParams(dataType, handle.GetDeviceName()); ss << " -DLOCAL_SIZE=" << std::to_string(wld); ss << " -DMIOPEN_USE_64BIT_INDEX=" << static_cast(use_64bit_index); @@ -1267,7 +1267,8 @@ void ScaleTensor(const Handle& handle, std::size_t wld = 256 < wgd ? 256 : wgd; // SUBTENSOR_OP_WITH_SCALAR set to 0 for set operation, and 1 for multiply operation - std::string parms = "-DSUBTENSOR_OP_WITH_SCALAR=1" + GetDataTypeKernelParams(dataType); + std::string parms = "-DSUBTENSOR_OP_WITH_SCALAR=1" + + GetDataTypeKernelParams(dataType, handle.GetDeviceName()); parms += " -DLOCAL_SIZE=" + std::to_string(wld); parms += " -DMIOPEN_USE_64BIT_INDEX=" + std::to_string(static_cast(use_64bit_index)); @@ -1432,7 +1433,8 @@ void CopyTensor(const Handle& handle, std::size_t wld = 256 < wgd ? 256 : wgd; - std::string parms = GetDataTypeKernelParams(srcDesc_flat.GetType()); + std::string parms = + GetDataTypeKernelParams(srcDesc_flat.GetType(), handle.GetDeviceName()); parms += " -DLOCAL_SIZE=" + std::to_string(wld); parms += @@ -2142,7 +2144,7 @@ void TransformTensor(const Handle& handle, std::size_t wld = 256 < wgd ? 256 : wgd; std::string parms = - GetDataTypeKernelParams(dataTypey) + + GetDataTypeKernelParams(dataTypey, handle.GetDeviceName()) + " -DMIOPEN_BETA_IS_ZERO=" + std::to_string(static_cast(is_beta_zero)) + " -DMIOPEN_ALPHA_IS_ONE=" + std::to_string(static_cast(is_alpha_one)); parms += " -DLOCAL_SIZE=" + std::to_string(wld); diff --git a/projects/miopen/test/gtest/solver_f8.hpp b/projects/miopen/test/gtest/solver_f8.hpp index b8b758bc989f..4371dcde72de 100644 --- a/projects/miopen/test/gtest/solver_f8.hpp +++ b/projects/miopen/test/gtest/solver_f8.hpp @@ -79,6 +79,8 @@ struct ConvFwdSolverTestF8 protected: void SetUp() override { + miopen_f8::set_hip_f8_ieee_exponent_bias( + miopen::UsesIeeeExponentBias(get_handle().GetDeviceName())); test_skipped = false; std::tie(algo, conv_config) = GetParam(); input = tensor{conv_config.N, conv_config.C, conv_config.H, conv_config.W}; From bc11c5d0bf70db92fbf708ad820099fa7f7c5c68 Mon Sep 17 00:00:00 2001 From: Yassin Soliman Date: Wed, 28 Oct 2026 17:35:27 -0600 Subject: [PATCH 3/4] test(miopen): set the host FP8 bias in the conv solver fixtures The FP8 conv solver tests generate FP8 data on the host and check the kernel output against a reference. Now that kernels get the exponent bias of the device they run on, the host values need the same bias, or gfx950 kernels read the host bits as values twice as large. Add SetHostFp8ExponentBias() to conv_test_base.hpp and call it from the forward, backward data and backward weights fixtures. Drop the call from ConvFwdSolverTestF8 in solver_f8.hpp, which no test instantiates. --- projects/miopen/test/gtest/conv_test_base.hpp | 15 +++++++++++++++ projects/miopen/test/gtest/solver_bwd.hpp | 1 + projects/miopen/test/gtest/solver_f8.hpp | 2 -- projects/miopen/test/gtest/solver_wrw.hpp | 1 + 4 files changed, 17 insertions(+), 2 deletions(-) diff --git a/projects/miopen/test/gtest/conv_test_base.hpp b/projects/miopen/test/gtest/conv_test_base.hpp index 224967c64a48..ada56d383ff2 100644 --- a/projects/miopen/test/gtest/conv_test_base.hpp +++ b/projects/miopen/test/gtest/conv_test_base.hpp @@ -27,14 +27,28 @@ #include #include +#include #include +#include #include "tensor_holder.hpp" #include "conv_common.hpp" #include "conv_tensor_gen.hpp" +#include "get_handle.hpp" using Direction = miopen::conv::Direction; +// Host FP8 values must use the exponent bias of the device that runs the kernels. +template +void SetHostFp8ExponentBias() +{ + if constexpr(std::is_same_v || std::is_same_v) + { + miopen_f8::set_hip_f8_ieee_exponent_bias( + miopen::UsesIeeeExponentBias(get_handle().GetDeviceName())); + } +} + struct GroupConvTestConfigBase { virtual ~GroupConvTestConfigBase() = default; @@ -377,6 +391,7 @@ struct ConvFwdSolverTestBase protected: void SetUpImpl(TConfig conv_config, miopenTensorLayout_t tensor_layout) { + SetHostFp8ExponentBias(); input = tensor{tensor_layout, conv_config.GetInput()}; weights = tensor{tensor_layout, conv_config.GetWeights()}; input.generate(GenData{}); diff --git a/projects/miopen/test/gtest/solver_bwd.hpp b/projects/miopen/test/gtest/solver_bwd.hpp index ef80199067b5..80bf0cbe4e8c 100644 --- a/projects/miopen/test/gtest/solver_bwd.hpp +++ b/projects/miopen/test/gtest/solver_bwd.hpp @@ -101,6 +101,7 @@ struct ConvBwdSolverTest GTEST_SKIP(); } + SetHostFp8ExponentBias(); input = tensor{conv_config.GetInput()}; weights = tensor{conv_config.GetWeights()}; weights.generate(GenWeights{}); diff --git a/projects/miopen/test/gtest/solver_f8.hpp b/projects/miopen/test/gtest/solver_f8.hpp index 4371dcde72de..b8b758bc989f 100644 --- a/projects/miopen/test/gtest/solver_f8.hpp +++ b/projects/miopen/test/gtest/solver_f8.hpp @@ -79,8 +79,6 @@ struct ConvFwdSolverTestF8 protected: void SetUp() override { - miopen_f8::set_hip_f8_ieee_exponent_bias( - miopen::UsesIeeeExponentBias(get_handle().GetDeviceName())); test_skipped = false; std::tie(algo, conv_config) = GetParam(); input = tensor{conv_config.N, conv_config.C, conv_config.H, conv_config.W}; diff --git a/projects/miopen/test/gtest/solver_wrw.hpp b/projects/miopen/test/gtest/solver_wrw.hpp index 75c3eb0ce90b..fca9413f2947 100644 --- a/projects/miopen/test/gtest/solver_wrw.hpp +++ b/projects/miopen/test/gtest/solver_wrw.hpp @@ -101,6 +101,7 @@ struct ConvWrwSolverTest GTEST_SKIP(); } + SetHostFp8ExponentBias(); input = tensor{conv_config.GetInput()}; weights = tensor{conv_config.GetWeights()}; input.generate(GenData{}); From e697ccf10e840b36589de4bb55509c45c37dc555 Mon Sep 17 00:00:00 2001 From: Yassin Soliman Date: Thu, 29 Oct 2026 10:32:19 -0600 Subject: [PATCH 4/4] [MIOpen] Add CPU tests for the per-device FP8 exponent bias Check that GetDataTypeKBP and GetDataTypeKernelParams pick bias 0 for gfx942 and 1 for gfx950 and gfx1200, and that SetHostFp8ExponentBias switches the host bias 0 to 1 to 0 in one process. SetHostFp8ExponentBias gets an overload that takes the device name so the test does not need a handle. --- projects/miopen/test/gtest/conv_test_base.hpp | 13 ++- .../test/gtest/unit_fp8_exponent_bias.cpp | 91 +++++++++++++++++++ 2 files changed, 102 insertions(+), 2 deletions(-) create mode 100644 projects/miopen/test/gtest/unit_fp8_exponent_bias.cpp diff --git a/projects/miopen/test/gtest/conv_test_base.hpp b/projects/miopen/test/gtest/conv_test_base.hpp index ada56d383ff2..828778a7a64c 100644 --- a/projects/miopen/test/gtest/conv_test_base.hpp +++ b/projects/miopen/test/gtest/conv_test_base.hpp @@ -29,6 +29,7 @@ #include #include #include +#include #include #include "tensor_holder.hpp" @@ -39,13 +40,21 @@ using Direction = miopen::conv::Direction; // Host FP8 values must use the exponent bias of the device that runs the kernels. +template +void SetHostFp8ExponentBias(const std::string& device_name) +{ + if constexpr(std::is_same_v || std::is_same_v) + { + miopen_f8::set_hip_f8_ieee_exponent_bias(miopen::UsesIeeeExponentBias(device_name)); + } +} + template void SetHostFp8ExponentBias() { if constexpr(std::is_same_v || std::is_same_v) { - miopen_f8::set_hip_f8_ieee_exponent_bias( - miopen::UsesIeeeExponentBias(get_handle().GetDeviceName())); + SetHostFp8ExponentBias(get_handle().GetDeviceName()); } } diff --git a/projects/miopen/test/gtest/unit_fp8_exponent_bias.cpp b/projects/miopen/test/gtest/unit_fp8_exponent_bias.cpp new file mode 100644 index 000000000000..309c58e3b63d --- /dev/null +++ b/projects/miopen/test/gtest/unit_fp8_exponent_bias.cpp @@ -0,0 +1,91 @@ +/******************************************************************************* + * + * MIT License + * + * Copyright (c) 2026 Advanced Micro Devices, Inc. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + * + *******************************************************************************/ +#include + +#include + +#include "conv_test_base.hpp" + +#include + +namespace { + +std::string BiasDefine(miopenDataType_t type, const std::string& device_name) +{ + return miopen::GetDataTypeKBP(type, device_name).GenerateFor(miopen::kbp::HIP{}); +} + +void ExpectKernelBias(const std::string& device_name, int expected_bias) +{ + const auto expected = "-DMIOPEN_FP8_IEEE_EXPONENT_BIAS=" + std::to_string(expected_bias); + for(const auto type : {miopenFloat8_fnuz, miopenBFloat8_fnuz}) + { + EXPECT_NE(BiasDefine(type, device_name).find(expected), std::string::npos) + << device_name << " type " << type << " expected " << expected; + EXPECT_NE(miopen::GetDataTypeKernelParams(type, device_name).find(expected), + std::string::npos) + << device_name << " type " << type << " expected " << expected; + } +} + +// Restores the process-wide host bias so other tests are unaffected. +struct HostBiasGuard +{ + bool saved = miopen_f8::get_hip_f8_ieee_exponent_bias(); + ~HostBiasGuard() { miopen_f8::set_hip_f8_ieee_exponent_bias(saved); } +}; + +} // namespace + +TEST(CPU_Fp8ExponentBias_NONE, KernelParamsUseDeviceBias) +{ + ExpectKernelBias("gfx942", 0); + ExpectKernelBias("gfx950", 1); + ExpectKernelBias("gfx1200", 1); +} + +TEST(CPU_Fp8ExponentBias_NONE, HostBiasSwitchesInOneProcess) +{ + HostBiasGuard guard; + + // 1.0f encodes as 0x40 with the FNUZ bias and as 0x38 (FP8) or 0x3C (BF8) with the IEEE bias. + const auto check = [](bool ieee) { + EXPECT_EQ(miopen_f8::get_hip_f8_ieee_exponent_bias(), ieee); + EXPECT_EQ(float8_fnuz(1.0f).data, ieee ? 0x38 : 0x40); + EXPECT_EQ(bfloat8_fnuz(1.0f).data, ieee ? 0x3C : 0x40); + }; + + SetHostFp8ExponentBias("gfx942"); + check(false); + SetHostFp8ExponentBias("gfx950"); + check(true); + SetHostFp8ExponentBias("gfx942"); + check(false); + SetHostFp8ExponentBias("gfx1200"); + check(true); + SetHostFp8ExponentBias("gfx942"); + check(false); +}