From d9b0d061725412951c47dce318dbf10d4823a297 Mon Sep 17 00:00:00 2001 From: Ilija Milenkovic Date: Fri, 25 Apr 2025 09:24:53 -0400 Subject: [PATCH 1/5] [release/2.6] Change preferred BLAS backend for gfx110x --- aten/src/ATen/Context.h | 17 ++++++++++++++--- 1 file changed, 14 insertions(+), 3 deletions(-) diff --git a/aten/src/ATen/Context.h b/aten/src/ATen/Context.h index 659c1a436e1eb..2d8c5fbf31d78 100644 --- a/aten/src/ATen/Context.h +++ b/aten/src/ATen/Context.h @@ -418,14 +418,25 @@ class TORCH_API Context { c10::utils::check_env("TORCH_LINALG_PREFER_CUSOLVER") == true ? at::LinalgBackend::Cusolver : at::LinalgBackend::Default; +#ifdef USE_ROCM + const bool is_gfx110x = []() { + const std::vector archs = {"gfx1100", "gfx1101"}; + for (auto index: c10::irange(detail::getCUDAHooks().getNumGPUs())) { + if (detail::getCUDAHooks().isGPUArch(index, archs)) { + return true; + } + } + return false; + }(); +#endif at::BlasBackend blas_preferred_backend = #ifdef USE_ROCM - (c10::utils::check_env("TORCH_BLAS_PREFER_HIPBLASLT") != false) + ((is_gfx110x && c10::utils::check_env("TORCH_BLAS_PREFER_HIPBLASLT") != true) || c10::utils::check_env("TORCH_BLAS_PREFER_HIPBLASLT") == false) #else (c10::utils::check_env("TORCH_BLAS_PREFER_CUBLASLT") == true) #endif - ? at::BlasBackend::Cublaslt - : at::BlasBackend::Cublas; + ? at::BlasBackend::Cublas + : at::BlasBackend::Cublaslt; at::ROCmFABackend rocm_fa_preferred_backend = c10::utils::check_env("TORCH_ROCM_FA_PREFER_CK") == true ? at::ROCmFABackend::Ck From 6efc0d8378a550603c71ba374c88b91c2298d62f Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Wed, 7 May 2025 16:56:54 +0000 Subject: [PATCH 2/5] Revert "[release/2.6] Change preferred BLAS backend for gfx110x" This reverts commit d9b0d061725412951c47dce318dbf10d4823a297. --- aten/src/ATen/Context.h | 17 +++-------------- 1 file changed, 3 insertions(+), 14 deletions(-) diff --git a/aten/src/ATen/Context.h b/aten/src/ATen/Context.h index 2d8c5fbf31d78..659c1a436e1eb 100644 --- a/aten/src/ATen/Context.h +++ b/aten/src/ATen/Context.h @@ -418,25 +418,14 @@ class TORCH_API Context { c10::utils::check_env("TORCH_LINALG_PREFER_CUSOLVER") == true ? at::LinalgBackend::Cusolver : at::LinalgBackend::Default; -#ifdef USE_ROCM - const bool is_gfx110x = []() { - const std::vector archs = {"gfx1100", "gfx1101"}; - for (auto index: c10::irange(detail::getCUDAHooks().getNumGPUs())) { - if (detail::getCUDAHooks().isGPUArch(index, archs)) { - return true; - } - } - return false; - }(); -#endif at::BlasBackend blas_preferred_backend = #ifdef USE_ROCM - ((is_gfx110x && c10::utils::check_env("TORCH_BLAS_PREFER_HIPBLASLT") != true) || c10::utils::check_env("TORCH_BLAS_PREFER_HIPBLASLT") == false) + (c10::utils::check_env("TORCH_BLAS_PREFER_HIPBLASLT") != false) #else (c10::utils::check_env("TORCH_BLAS_PREFER_CUBLASLT") == true) #endif - ? at::BlasBackend::Cublas - : at::BlasBackend::Cublaslt; + ? at::BlasBackend::Cublaslt + : at::BlasBackend::Cublas; at::ROCmFABackend rocm_fa_preferred_backend = c10::utils::check_env("TORCH_ROCM_FA_PREFER_CK") == true ? at::ROCmFABackend::Ck From 2a1375e50893f332fb2a3f5d74556b7359032f42 Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Wed, 7 May 2025 16:57:13 +0000 Subject: [PATCH 3/5] next attempt --- aten/src/ATen/Context.h | 32 +++++++++++++++++++++++++++++--- 1 file changed, 29 insertions(+), 3 deletions(-) diff --git a/aten/src/ATen/Context.h b/aten/src/ATen/Context.h index 659c1a436e1eb..4411b233a2be2 100644 --- a/aten/src/ATen/Context.h +++ b/aten/src/ATen/Context.h @@ -418,12 +418,38 @@ class TORCH_API Context { c10::utils::check_env("TORCH_LINALG_PREFER_CUSOLVER") == true ? at::LinalgBackend::Cusolver : at::LinalgBackend::Default; - at::BlasBackend blas_preferred_backend = #ifdef USE_ROCM - (c10::utils::check_env("TORCH_BLAS_PREFER_HIPBLASLT") != false) + // AMD Instinct targets prefer hipblaslt + static const bool _hipblaslt_preferred_default = []() { + static const std::vector archs = { + "gfx90a", "gfx942", +#if ROCM_VERSION >= 60500 + "gfx950" +#endif + }; + for (auto index: c10::irange(detail::getCUDAHooks().deviceCount())) { + if (!detail::getCUDAHooks().isGPUArch(archs, index)) { + return false; + } + } + return true; + }(); #else - (c10::utils::check_env("TORCH_BLAS_PREFER_CUBLASLT") == true) + static const bool _hipblaslt_preferred_default = false; +#endif + static const bool _blaslt_preferred = []() { + if (c10::utils::check_env("TORCH_BLAS_PREFER_CUBLASLT") == true) { + return true; + } + if (c10::utils::check_env("TORCH_BLAS_PREFER_HIPBLASLT") == true) { + return true; + } +#ifdef USE_ROCM + return _hipblaslt_preferred_default; #endif + return false; + }(); + at::BlasBackend blas_preferred_backend = _blaslt_preferred ? at::BlasBackend::Cublaslt : at::BlasBackend::Cublas; at::ROCmFABackend rocm_fa_preferred_backend = From 744671327dcced25d3f72ab8bc7c86e0385106eb Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Wed, 7 May 2025 17:01:30 +0000 Subject: [PATCH 4/5] allow user override to false --- aten/src/ATen/Context.h | 10 ++++++---- 1 file changed, 6 insertions(+), 4 deletions(-) diff --git a/aten/src/ATen/Context.h b/aten/src/ATen/Context.h index 4411b233a2be2..d7e7c77db5838 100644 --- a/aten/src/ATen/Context.h +++ b/aten/src/ATen/Context.h @@ -438,11 +438,13 @@ class TORCH_API Context { static const bool _hipblaslt_preferred_default = false; #endif static const bool _blaslt_preferred = []() { - if (c10::utils::check_env("TORCH_BLAS_PREFER_CUBLASLT") == true) { - return true; + auto env = c10::utils::check_env("TORCH_BLAS_PREFER_CUBLASLT"); + if (env.has_value()) { + return env.value(); } - if (c10::utils::check_env("TORCH_BLAS_PREFER_HIPBLASLT") == true) { - return true; + env = c10::utils::check_env("TORCH_BLAS_PREFER_HIPBLASLT"); + if (env.has_value()) { + return env.value(); } #ifdef USE_ROCM return _hipblaslt_preferred_default; From 14341d582f9184f6b9556e4252bbe2ccd921e3c6 Mon Sep 17 00:00:00 2001 From: Ilija Milenkovic Date: Mon, 12 May 2025 10:46:52 -0400 Subject: [PATCH 5/5] Add gfx12xx to _hipblaslt_preferred_default list and apply fixes --- aten/src/ATen/Context.h | 13 ++++++++----- 1 file changed, 8 insertions(+), 5 deletions(-) diff --git a/aten/src/ATen/Context.h b/aten/src/ATen/Context.h index d7e7c77db5838..444dfd1b02160 100644 --- a/aten/src/ATen/Context.h +++ b/aten/src/ATen/Context.h @@ -420,24 +420,27 @@ class TORCH_API Context { : at::LinalgBackend::Default; #ifdef USE_ROCM // AMD Instinct targets prefer hipblaslt - static const bool _hipblaslt_preferred_default = []() { - static const std::vector archs = { + const bool _hipblaslt_preferred_default = []() { + const std::vector archs = { "gfx90a", "gfx942", +#if ROCM_VERSION >= 60300 + "gfx1200", "gfx1201", +#endif #if ROCM_VERSION >= 60500 "gfx950" #endif }; for (auto index: c10::irange(detail::getCUDAHooks().deviceCount())) { - if (!detail::getCUDAHooks().isGPUArch(archs, index)) { + if (!detail::getCUDAHooks().isGPUArch(index, archs)) { return false; } } return true; }(); #else - static const bool _hipblaslt_preferred_default = false; + const bool _hipblaslt_preferred_default = false; #endif - static const bool _blaslt_preferred = []() { + const bool _blaslt_preferred = [&]() { auto env = c10::utils::check_env("TORCH_BLAS_PREFER_CUBLASLT"); if (env.has_value()) { return env.value();