diff --git a/projects/miopen/CHANGELOG.md b/projects/miopen/CHANGELOG.md index 971e66fbcbbf..d9033c3cd903 100644 --- a/projects/miopen/CHANGELOG.md +++ b/projects/miopen/CHANGELOG.md @@ -30,6 +30,10 @@ Full documentation for MIOpen is available [here](https://rocm.docs.amd.com/proj the 2 GiB bound are restored where the tuned config passes the new check. This restores ResNet-50/101 FP32 batch-1024 backward-weights speed. FP16/BF16 keep the 2 GiB bound, and so does the weight tensor in every case. +* [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. * [Conv] Fixed HIP graph captures of convolutions failing with `hipErrorStreamCaptureInvalidated` after a find had skipped a naive solver for exceeding `MIOPEN_NAIVE_TIMEOUT` (#12121). Scratch held by a skipped evaluation is now released by a background thread as soon as its kernel @@ -55,6 +59,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. ### Optimized * [Conv] Restored gfx1200 backward-weights performance that #13697 gave up when it removed the 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 3d1b42b147d2..776ce9f3a6bf 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/conv_test_base.hpp b/projects/miopen/test/gtest/conv_test_base.hpp index 224967c64a48..828778a7a64c 100644 --- a/projects/miopen/test/gtest/conv_test_base.hpp +++ b/projects/miopen/test/gtest/conv_test_base.hpp @@ -27,14 +27,37 @@ #include #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(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) + { + SetHostFp8ExponentBias(get_handle().GetDeviceName()); + } +} + struct GroupConvTestConfigBase { virtual ~GroupConvTestConfigBase() = default; @@ -377,6 +400,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/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 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_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{}); 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); +}