diff --git a/CMakeLists.txt b/CMakeLists.txt index 8291f5a3a6b..dfc7421ceca 100755 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -323,7 +323,6 @@ rocm_enable_cppcheck( ${CMAKE_CURRENT_SOURCE_DIR}/test/include DEFINE MIGRAPHX_MLIR=1 - MIGRAPHX_ENABLE_HIPBLASLT_GEMM=1 MIGRAPHX_HAS_EXECUTORS=0 CPPCHECK=1 MIGRAPHX_USE_MIOPEN=1 diff --git a/docs/dev/env_vars.rst b/docs/dev/env_vars.rst index dd2fd7b4efd..0bc73cbb217 100644 --- a/docs/dev/env_vars.rst +++ b/docs/dev/env_vars.rst @@ -135,8 +135,8 @@ Use it in conjunction with ``MIGRAPHX_DISABLE_MLIR=1``. Set to "1", "enable", "enabled", "yes", or "true" to use. Disables use of the rocMLIR library. -.. envvar:: MIGRAPHX_ENABLE_HIPBLASLT_GEMM -Set to "0", "disable", "disabled", "no", or "false" to use. +.. envvar:: MIGRAPHX_DISABLE_HIPBLASLT_GEMM +Set to "1", "enable", "enabled", "yes", or "true" to use. Disables use of hipBLASLt. .. envvar:: MIGRAPHX_COPY_LITERALS diff --git a/src/targets/gpu/fuse_ops.cpp b/src/targets/gpu/fuse_ops.cpp index 5e93ccf5ecf..25c5e562c3b 100644 --- a/src/targets/gpu/fuse_ops.cpp +++ b/src/targets/gpu/fuse_ops.cpp @@ -43,7 +43,6 @@ namespace migraphx { inline namespace MIGRAPHX_INLINE_NS { namespace gpu { -MIGRAPHX_DECLARE_ENV_VAR(MIGRAPHX_ENABLE_HIPBLASLT_GEMM) MIGRAPHX_DECLARE_ENV_VAR(MIGRAPHX_DISABLE_MIOPEN_FUSION) #if MIGRAPHX_USE_MIOPEN struct fusion diff --git a/src/targets/gpu/lowering.cpp b/src/targets/gpu/lowering.cpp index 58820240681..7390f88abe5 100644 --- a/src/targets/gpu/lowering.cpp +++ b/src/targets/gpu/lowering.cpp @@ -55,7 +55,7 @@ namespace migraphx { inline namespace MIGRAPHX_INLINE_NS { namespace gpu { -MIGRAPHX_DECLARE_ENV_VAR(MIGRAPHX_ENABLE_HIPBLASLT_GEMM); +MIGRAPHX_DECLARE_ENV_VAR(MIGRAPHX_DISABLE_HIPBLASLT_GEMM); MIGRAPHX_DECLARE_ENV_VAR(MIGRAPHX_DISABLE_MIOPEN_POOLING) struct miopen_apply @@ -254,7 +254,7 @@ struct miopen_apply auto output = insert_allocation(ins, ins->get_shape()); refs.push_back(output); #if MIGRAPHX_USE_HIPBLASLT - if(disabled(MIGRAPHX_ENABLE_HIPBLASLT_GEMM{}) or not hipblaslt_supported()) + if(enabled(MIGRAPHX_DISABLE_HIPBLASLT_GEMM{}) or not hipblaslt_supported()) { #endif return mod->replace_instruction( diff --git a/src/targets/gpu/target.cpp b/src/targets/gpu/target.cpp index bdafff1877e..d937dfaf54f 100644 --- a/src/targets/gpu/target.cpp +++ b/src/targets/gpu/target.cpp @@ -82,7 +82,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) +MIGRAPHX_DECLARE_ENV_VAR(MIGRAPHX_DISABLE_HIPBLASLT_GEMM) std::vector target::get_passes(migraphx::context& gctx, const compile_options& options) const { @@ -107,7 +107,7 @@ std::vector target::get_passes(migraphx::context& gctx, const compile_opti // different between fp8e4m3fnuz and OCP types because rocBLAS only has // support for fp8e4m3fnuz std::set unsupported_fp8e4m3fnuz_ops = {}; - if(disabled(MIGRAPHX_ENABLE_HIPBLASLT_GEMM{}) and not gpu::rocblas_fp8_available()) + if(enabled(MIGRAPHX_DISABLE_HIPBLASLT_GEMM{}) and not gpu::rocblas_fp8_available()) { unsupported_fp8e4m3fnuz_ops.insert("dot"); unsupported_fp8e4m3fnuz_ops.insert("quant_dot"); @@ -134,19 +134,14 @@ std::vector target::get_passes(migraphx::context& gctx, const compile_opti std::set unsupported_fp8e5m2fnuz_ops = unsupported_fp8e4m3fnuz_ops; // disable gemm for fp8e5m2fnuz if rocBLAS is being used - if(disabled(MIGRAPHX_ENABLE_HIPBLASLT_GEMM{})) + if(enabled(MIGRAPHX_DISABLE_HIPBLASLT_GEMM{})) { unsupported_fp8e5m2fnuz_ops.insert("dot"); unsupported_fp8e5m2fnuz_ops.insert("quant_dot"); } std::set unsupported_fp8ocp_ops = {}; - // TODO: remove this when the flag is removed - if(disabled(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");