From e6e3a5a7f6be72f4fd17fb8e91f8e9aeb342de0b Mon Sep 17 00:00:00 2001 From: charlie Date: Thu, 3 Oct 2024 11:00:41 -0500 Subject: [PATCH 01/12] Initial --- src/targets/gpu/hip_gemm_impl.cpp | 4 ++-- src/targets/gpu/target.cpp | 8 ++++++-- 2 files changed, 8 insertions(+), 4 deletions(-) diff --git a/src/targets/gpu/hip_gemm_impl.cpp b/src/targets/gpu/hip_gemm_impl.cpp index f5ec898d8d5..2f593da70ae 100644 --- a/src/targets/gpu/hip_gemm_impl.cpp +++ b/src/targets/gpu/hip_gemm_impl.cpp @@ -70,8 +70,8 @@ hipDataType get_type_hipblas(shape::type_t type) case shape::int32_type: return HIP_R_32I; case shape::uint32_type: return HIP_R_32U; case shape::fp8e4m3fnuz_type: return HIP_R_8F_E4M3_FNUZ; - case shape::fp8e4m3fn_type: - case shape::fp8e5m2_type: + case shape::fp8e4m3fn_type: return HIP_R_8F_E4M3; + case shape::fp8e5m2_type: return HIP_R_8F_E5M2; case shape::tuple_type: case shape::bool_type: case shape::uint16_type: diff --git a/src/targets/gpu/target.cpp b/src/targets/gpu/target.cpp index b190d11a402..8707e333f9a 100644 --- a/src/targets/gpu/target.cpp +++ b/src/targets/gpu/target.cpp @@ -130,8 +130,12 @@ std::vector target::get_passes(migraphx::context& gctx, const compile_opti std::set unsupported_fp8ocp_ops = {}; // TODO update with hipBLASLt support - unsupported_fp8ocp_ops.insert("dot"); - unsupported_fp8ocp_ops.insert("quant_dot"); +#if !MIGRAPHX_ENABLE_HIPBLASLT_GEMM + { + unsupported_fp8ocp_ops.insert("dot"); + unsupported_fp8ocp_ops.insert("quant_dot"); + } +#endif #if MIGRAPHX_USE_MIOPEN // MIOpen doesn't have support for fp8 pooling yet. unsupported_fp8ocp_ops.insert("pooling"); From dd60b3957f413d952d85af793e84032caf87e9b9 Mon Sep 17 00:00:00 2001 From: charlie Date: Thu, 3 Oct 2024 13:32:48 -0500 Subject: [PATCH 02/12] Remove conditional for now --- src/targets/gpu/target.cpp | 7 ------- 1 file changed, 7 deletions(-) diff --git a/src/targets/gpu/target.cpp b/src/targets/gpu/target.cpp index 8707e333f9a..b4cb1126dc5 100644 --- a/src/targets/gpu/target.cpp +++ b/src/targets/gpu/target.cpp @@ -129,13 +129,6 @@ std::vector target::get_passes(migraphx::context& gctx, const compile_opti unsupported_fp8e4m3fnuz_ops.insert("argmin"); std::set unsupported_fp8ocp_ops = {}; - // TODO update with hipBLASLt support -#if !MIGRAPHX_ENABLE_HIPBLASLT_GEMM - { - unsupported_fp8ocp_ops.insert("dot"); - unsupported_fp8ocp_ops.insert("quant_dot"); - } -#endif #if MIGRAPHX_USE_MIOPEN // MIOpen doesn't have support for fp8 pooling yet. unsupported_fp8ocp_ops.insert("pooling"); From 3628c7ac4bf4b06ec88ad717441bcb7d125e5855 Mon Sep 17 00:00:00 2001 From: charlie Date: Wed, 30 Oct 2024 15:29:36 -0500 Subject: [PATCH 03/12] Add check for the hipblas gemm flag --- src/targets/gpu/target.cpp | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/src/targets/gpu/target.cpp b/src/targets/gpu/target.cpp index 35a28217f55..5cca7619214 100644 --- a/src/targets/gpu/target.cpp +++ b/src/targets/gpu/target.cpp @@ -129,6 +129,11 @@ std::vector target::get_passes(migraphx::context& gctx, const compile_opti unsupported_fp8e4m3fnuz_ops.insert("argmin"); std::set unsupported_fp8ocp_ops = {}; + // TODO: remove this when the flag is removed +#if !MIGRAPHX_ENABLE_HIPBLASLT_GEMM + unsupported_fp8ocp_ops.insert("dot"); + unsupported_fp8ocp_ops.insert("quant_dot"); +#endif #if MIGRAPHX_USE_MIOPEN // MIOpen doesn't have support for fp8 pooling yet. unsupported_fp8ocp_ops.insert("pooling"); From f28d40d397769f020c45f18dbd8b0eedb47768be Mon Sep 17 00:00:00 2001 From: charlie Date: Wed, 30 Oct 2024 16:13:35 -0500 Subject: [PATCH 04/12] Declare env var, disable fp8ocp gemm with no intrinsics --- src/targets/gpu/target.cpp | 3 +++ 1 file changed, 3 insertions(+) diff --git a/src/targets/gpu/target.cpp b/src/targets/gpu/target.cpp index 5cca7619214..b01f59db3b8 100644 --- a/src/targets/gpu/target.cpp +++ b/src/targets/gpu/target.cpp @@ -81,6 +81,7 @@ MIGRAPHX_DECLARE_ENV_VAR(MIGRAPHX_ENABLE_NHWC) #ifndef _WIN32 MIGRAPHX_DECLARE_ENV_VAR(MIGRAPHX_ENABLE_CK) #endif +MIGRAPHX_DECLARE_ENV_VAR(MIGRAPHX_ENABLE_HIPBLASLT_GEMM) std::vector target::get_passes(migraphx::context& gctx, const compile_options& options) const { @@ -142,6 +143,8 @@ std::vector target::get_passes(migraphx::context& gctx, const compile_opti { unsupported_fp8ocp_ops.insert("convolution"); unsupported_fp8ocp_ops.insert("quant_convolution"); + unsupported_fp8ocp_ops.insert("dot"); + unsupported_fp8ocp_ops.insert("quant_dot"); } // add all device kernels unsupported_fp8ocp_ops.insert("logsoftmax"); From 8cc9914b0d65f595f92e561c3c402d0adfe94fbd Mon Sep 17 00:00:00 2001 From: charlie Date: Fri, 1 Nov 2024 11:25:19 -0500 Subject: [PATCH 05/12] Use enabled --- src/targets/gpu/target.cpp | 9 +++++---- 1 file changed, 5 insertions(+), 4 deletions(-) diff --git a/src/targets/gpu/target.cpp b/src/targets/gpu/target.cpp index b01f59db3b8..e91ae02e929 100644 --- a/src/targets/gpu/target.cpp +++ b/src/targets/gpu/target.cpp @@ -131,10 +131,11 @@ std::vector target::get_passes(migraphx::context& gctx, const compile_opti std::set unsupported_fp8ocp_ops = {}; // TODO: remove this when the flag is removed -#if !MIGRAPHX_ENABLE_HIPBLASLT_GEMM - unsupported_fp8ocp_ops.insert("dot"); - unsupported_fp8ocp_ops.insert("quant_dot"); -#endif + if(not enabled(MIGRAPHX_ENABLE_HIPBLASLT_GEMM{})) + { + unsupported_fp8ocp_ops.insert("dot"); + unsupported_fp8ocp_ops.insert("quant_dot"); + } #if MIGRAPHX_USE_MIOPEN // MIOpen doesn't have support for fp8 pooling yet. unsupported_fp8ocp_ops.insert("pooling"); From 856b6cbeb9e4e1f719e3251d12bb4416de77bd71 Mon Sep 17 00:00:00 2001 From: charlie Date: Tue, 5 Nov 2024 13:20:09 -0600 Subject: [PATCH 06/12] Turn off unsupported hipblas types --- src/targets/gpu/hip_gemm_impl.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/targets/gpu/hip_gemm_impl.cpp b/src/targets/gpu/hip_gemm_impl.cpp index 2f593da70ae..f5ec898d8d5 100644 --- a/src/targets/gpu/hip_gemm_impl.cpp +++ b/src/targets/gpu/hip_gemm_impl.cpp @@ -70,8 +70,8 @@ hipDataType get_type_hipblas(shape::type_t type) case shape::int32_type: return HIP_R_32I; case shape::uint32_type: return HIP_R_32U; case shape::fp8e4m3fnuz_type: return HIP_R_8F_E4M3_FNUZ; - case shape::fp8e4m3fn_type: return HIP_R_8F_E4M3; - case shape::fp8e5m2_type: return HIP_R_8F_E5M2; + case shape::fp8e4m3fn_type: + case shape::fp8e5m2_type: case shape::tuple_type: case shape::bool_type: case shape::uint16_type: From 8e00c912b4cdfc6e91eff3efaf2eb66b435f8b7b Mon Sep 17 00:00:00 2001 From: charlie Date: Tue, 5 Nov 2024 13:26:12 -0600 Subject: [PATCH 07/12] Hipblaslt types check gfx number --- src/targets/gpu/hip_gemm_impl.cpp | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/src/targets/gpu/hip_gemm_impl.cpp b/src/targets/gpu/hip_gemm_impl.cpp index f5ec898d8d5..6e1ecf7e1d9 100644 --- a/src/targets/gpu/hip_gemm_impl.cpp +++ b/src/targets/gpu/hip_gemm_impl.cpp @@ -71,7 +71,15 @@ hipDataType get_type_hipblas(shape::type_t type) case shape::uint32_type: return HIP_R_32U; case shape::fp8e4m3fnuz_type: return HIP_R_8F_E4M3_FNUZ; case shape::fp8e4m3fn_type: + if(gfx_has_fp8ocp_intrinsics()) + { + return HIP_R_8F_E4M3; + } case shape::fp8e5m2_type: + if(gfx_has_fp8ocp_intrinsics()) + { + return HIP_R_8F_E5M2; + } case shape::tuple_type: case shape::bool_type: case shape::uint16_type: From af9d6e33278f69978c42b3bc6fdee522fa5e6760 Mon Sep 17 00:00:00 2001 From: charlie Date: Tue, 5 Nov 2024 13:39:36 -0600 Subject: [PATCH 08/12] revert to previous --- src/targets/gpu/hip_gemm_impl.cpp | 12 ++---------- 1 file changed, 2 insertions(+), 10 deletions(-) diff --git a/src/targets/gpu/hip_gemm_impl.cpp b/src/targets/gpu/hip_gemm_impl.cpp index 6e1ecf7e1d9..2f593da70ae 100644 --- a/src/targets/gpu/hip_gemm_impl.cpp +++ b/src/targets/gpu/hip_gemm_impl.cpp @@ -70,16 +70,8 @@ hipDataType get_type_hipblas(shape::type_t type) case shape::int32_type: return HIP_R_32I; case shape::uint32_type: return HIP_R_32U; case shape::fp8e4m3fnuz_type: return HIP_R_8F_E4M3_FNUZ; - case shape::fp8e4m3fn_type: - if(gfx_has_fp8ocp_intrinsics()) - { - return HIP_R_8F_E4M3; - } - case shape::fp8e5m2_type: - if(gfx_has_fp8ocp_intrinsics()) - { - return HIP_R_8F_E5M2; - } + case shape::fp8e4m3fn_type: return HIP_R_8F_E4M3; + case shape::fp8e5m2_type: return HIP_R_8F_E5M2; case shape::tuple_type: case shape::bool_type: case shape::uint16_type: From 223ba4d6b6894c456e005a5a8c0b10f19180285d Mon Sep 17 00:00:00 2001 From: charlie Date: Tue, 5 Nov 2024 17:05:42 -0600 Subject: [PATCH 09/12] add preprocessor conditional on hipblaslt version --- src/targets/gpu/hip_gemm_impl.cpp | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/src/targets/gpu/hip_gemm_impl.cpp b/src/targets/gpu/hip_gemm_impl.cpp index 2f593da70ae..ea85b2452ce 100644 --- a/src/targets/gpu/hip_gemm_impl.cpp +++ b/src/targets/gpu/hip_gemm_impl.cpp @@ -70,8 +70,15 @@ hipDataType get_type_hipblas(shape::type_t type) case shape::int32_type: return HIP_R_32I; case shape::uint32_type: return HIP_R_32U; case shape::fp8e4m3fnuz_type: return HIP_R_8F_E4M3_FNUZ; +// TODO remove this preprocessor conditional when hipblaslt verison defaults to > 0.10.0 +#if(HIPBLASLT_VERSION_MAJOR * 100000 + HIPBLASLT_VERSION_MINOR * 100 + HIPBLASLT_VERSION_PATCH) < \ + 1000 + case shape::fp8e4m3fn_type: + case shape::fp8e5m2_type: +#else case shape::fp8e4m3fn_type: return HIP_R_8F_E4M3; case shape::fp8e5m2_type: return HIP_R_8F_E5M2; +#endif case shape::tuple_type: case shape::bool_type: case shape::uint16_type: From 7f5798c4f325d5fbf42c8cf722a5abaa4542d063 Mon Sep 17 00:00:00 2001 From: charlie Date: Thu, 7 Nov 2024 13:59:02 -0600 Subject: [PATCH 10/12] Use CMake macro to check symbols (needs testing) --- src/targets/gpu/CMakeLists.txt | 5 ++++- src/targets/gpu/hip_gemm_impl.cpp | 9 ++++----- 2 files changed, 8 insertions(+), 6 deletions(-) diff --git a/src/targets/gpu/CMakeLists.txt b/src/targets/gpu/CMakeLists.txt index 756fc909035..3e377698e26 100644 --- a/src/targets/gpu/CMakeLists.txt +++ b/src/targets/gpu/CMakeLists.txt @@ -22,6 +22,8 @@ # THE SOFTWARE. # #################################################################################### +include(CheckSymbolExists) + find_package(hip REQUIRED) if(NOT GPU_TARGETS) set(fatal_msg "HIP package is broken and has no GPU_TARGETS. Please pass GPU_TARGETS to cmake.") @@ -52,7 +54,8 @@ if(MIGRAPHX_USE_HIPBLASLT) # Making hipblas required to workaround the broken hipblaslt package. find_package(hipblas REQUIRED) message(STATUS "MIGraphx build with hipBLAS and hipBLASLt") -else() + check_symbol_exists(HIP_R_8F_E4M3 "hipblaslt.h" MIGRAPHX_HIPBLASLT_HAS_8F_E4M3) + check_symbol_exists(HIP_R_8F_E5M2 "hipblaslt.h" MIGRAPHX_HIPBLASLT_HAS_8F_E5M2) message(STATUS "MIGraphX build without hipBLAS and hipBLASLt") endif() diff --git a/src/targets/gpu/hip_gemm_impl.cpp b/src/targets/gpu/hip_gemm_impl.cpp index ea85b2452ce..dd90d1745bc 100644 --- a/src/targets/gpu/hip_gemm_impl.cpp +++ b/src/targets/gpu/hip_gemm_impl.cpp @@ -71,13 +71,12 @@ hipDataType get_type_hipblas(shape::type_t type) case shape::uint32_type: return HIP_R_32U; case shape::fp8e4m3fnuz_type: return HIP_R_8F_E4M3_FNUZ; // TODO remove this preprocessor conditional when hipblaslt verison defaults to > 0.10.0 -#if(HIPBLASLT_VERSION_MAJOR * 100000 + HIPBLASLT_VERSION_MINOR * 100 + HIPBLASLT_VERSION_PATCH) < \ - 1000 - case shape::fp8e4m3fn_type: - case shape::fp8e5m2_type: -#else +#if defined(MIGRAPHX_HIPBLASLT_HAS_8F_E4M3) && defined(MIGRAPHX_HIPBLASLT_HAS_8F_E5M2) case shape::fp8e4m3fn_type: return HIP_R_8F_E4M3; case shape::fp8e5m2_type: return HIP_R_8F_E5M2; +#else + case shape::fp8e4m3fn_type: + case shape::fp8e5m2_type: #endif case shape::tuple_type: case shape::bool_type: From 0506abdedccd2ea52e2cbfa94179f9a123ff6608 Mon Sep 17 00:00:00 2001 From: charlie Date: Thu, 7 Nov 2024 15:05:39 -0600 Subject: [PATCH 11/12] Revert " Use CMake macro to check symbols (needs testing)" This reverts commit 7f5798c4f325d5fbf42c8cf722a5abaa4542d063. --- src/targets/gpu/CMakeLists.txt | 5 +---- src/targets/gpu/hip_gemm_impl.cpp | 9 +++++---- 2 files changed, 6 insertions(+), 8 deletions(-) diff --git a/src/targets/gpu/CMakeLists.txt b/src/targets/gpu/CMakeLists.txt index 3e377698e26..756fc909035 100644 --- a/src/targets/gpu/CMakeLists.txt +++ b/src/targets/gpu/CMakeLists.txt @@ -22,8 +22,6 @@ # THE SOFTWARE. # #################################################################################### -include(CheckSymbolExists) - find_package(hip REQUIRED) if(NOT GPU_TARGETS) set(fatal_msg "HIP package is broken and has no GPU_TARGETS. Please pass GPU_TARGETS to cmake.") @@ -54,8 +52,7 @@ if(MIGRAPHX_USE_HIPBLASLT) # Making hipblas required to workaround the broken hipblaslt package. find_package(hipblas REQUIRED) message(STATUS "MIGraphx build with hipBLAS and hipBLASLt") - check_symbol_exists(HIP_R_8F_E4M3 "hipblaslt.h" MIGRAPHX_HIPBLASLT_HAS_8F_E4M3) - check_symbol_exists(HIP_R_8F_E5M2 "hipblaslt.h" MIGRAPHX_HIPBLASLT_HAS_8F_E5M2) +else() message(STATUS "MIGraphX build without hipBLAS and hipBLASLt") endif() diff --git a/src/targets/gpu/hip_gemm_impl.cpp b/src/targets/gpu/hip_gemm_impl.cpp index dd90d1745bc..ea85b2452ce 100644 --- a/src/targets/gpu/hip_gemm_impl.cpp +++ b/src/targets/gpu/hip_gemm_impl.cpp @@ -71,12 +71,13 @@ hipDataType get_type_hipblas(shape::type_t type) case shape::uint32_type: return HIP_R_32U; case shape::fp8e4m3fnuz_type: return HIP_R_8F_E4M3_FNUZ; // TODO remove this preprocessor conditional when hipblaslt verison defaults to > 0.10.0 -#if defined(MIGRAPHX_HIPBLASLT_HAS_8F_E4M3) && defined(MIGRAPHX_HIPBLASLT_HAS_8F_E5M2) - case shape::fp8e4m3fn_type: return HIP_R_8F_E4M3; - case shape::fp8e5m2_type: return HIP_R_8F_E5M2; -#else +#if(HIPBLASLT_VERSION_MAJOR * 100000 + HIPBLASLT_VERSION_MINOR * 100 + HIPBLASLT_VERSION_PATCH) < \ + 1000 case shape::fp8e4m3fn_type: case shape::fp8e5m2_type: +#else + case shape::fp8e4m3fn_type: return HIP_R_8F_E4M3; + case shape::fp8e5m2_type: return HIP_R_8F_E5M2; #endif case shape::tuple_type: case shape::bool_type: From d453fdc3f389dafedc1a769abe39d5e8cd8d401c Mon Sep 17 00:00:00 2001 From: charlie Date: Fri, 8 Nov 2024 18:00:00 -0600 Subject: [PATCH 12/12] Use variable defined by hipblaslt for HIP float8 ocp types --- src/targets/gpu/hip_gemm_impl.cpp | 11 +++++------ 1 file changed, 5 insertions(+), 6 deletions(-) diff --git a/src/targets/gpu/hip_gemm_impl.cpp b/src/targets/gpu/hip_gemm_impl.cpp index ea85b2452ce..4e282cc01ff 100644 --- a/src/targets/gpu/hip_gemm_impl.cpp +++ b/src/targets/gpu/hip_gemm_impl.cpp @@ -70,14 +70,13 @@ hipDataType get_type_hipblas(shape::type_t type) case shape::int32_type: return HIP_R_32I; case shape::uint32_type: return HIP_R_32U; case shape::fp8e4m3fnuz_type: return HIP_R_8F_E4M3_FNUZ; -// TODO remove this preprocessor conditional when hipblaslt verison defaults to > 0.10.0 -#if(HIPBLASLT_VERSION_MAJOR * 100000 + HIPBLASLT_VERSION_MINOR * 100 + HIPBLASLT_VERSION_PATCH) < \ - 1000 - case shape::fp8e4m3fn_type: - case shape::fp8e5m2_type: -#else +// TODO can remove this preprocessor conditional when hip verison defaults to have these types +#ifdef ROCM_USE_FLOAT8 case shape::fp8e4m3fn_type: return HIP_R_8F_E4M3; case shape::fp8e5m2_type: return HIP_R_8F_E5M2; +#else + case shape::fp8e4m3fn_type: + case shape::fp8e5m2_type: #endif case shape::tuple_type: case shape::bool_type: