From 060691eb81100616f5403fdfdc1826096b796aba Mon Sep 17 00:00:00 2001
From: Ilija Milenkovic <imilenko@amd.com>
Date: Tue, 13 May 2025 11:44:23 +0200
Subject: [PATCH 1/2] [release/2.6] Change gfx110x BLAS preferred backend
 (#2053)

Only AMD Instinct GPUs and Navi 4x prefer hipblaslt by default, but user can still
override using env var.

---------

Co-authored-by: Jeff Daily <jeff.daily@amd.com>
---
 aten/src/ATen/Context.h | 37 ++++++++++++++++++++++++++++++++++---
 1 file changed, 34 insertions(+), 3 deletions(-)

diff --git a/aten/src/ATen/Context.h b/aten/src/ATen/Context.h
index d3ffb51d49659..911d4a9688e2c 100644
--- a/aten/src/ATen/Context.h
+++ b/aten/src/ATen/Context.h
@@ -425,12 +425,43 @@ 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
+  const bool _hipblaslt_preferred_default = []() {
+    const std::vector<std::string> 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(index, archs)) {
+        return false;
+      }
+    }
+    return true;
+  }();
 #else
-      (c10::utils::check_env("TORCH_BLAS_PREFER_CUBLASLT") == true)
+  const bool _hipblaslt_preferred_default = false;
+#endif
+  const bool _blaslt_preferred = [&]() {
+    auto env = c10::utils::check_env("TORCH_BLAS_PREFER_CUBLASLT");
+    if (env.has_value()) {
+      return env.value();
+    }
+    env = c10::utils::check_env("TORCH_BLAS_PREFER_HIPBLASLT");
+    if (env.has_value()) {
+      return env.value();
+    }
+#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 6b3eece6dd26c135287dba4f210e72a763f5c4c7 Mon Sep 17 00:00:00 2001
From: Arash Pakbin <Arash.Pakbin@amd.com>
Date: Wed, 21 May 2025 18:01:34 +0000
Subject: [PATCH 2/2] defaulting to hipBLASLt for gfx90a and gfx942 only

---
 aten/src/ATen/Context.h | 8 +-------
 1 file changed, 1 insertion(+), 7 deletions(-)

diff --git a/aten/src/ATen/Context.h b/aten/src/ATen/Context.h
index 911d4a9688e2c..5060bf7f6aabd 100644
--- a/aten/src/ATen/Context.h
+++ b/aten/src/ATen/Context.h
@@ -429,13 +429,7 @@ class TORCH_API Context {
   // AMD Instinct targets prefer hipblaslt
   const bool _hipblaslt_preferred_default = []() {
     const std::vector<std::string> archs = {
-        "gfx90a", "gfx942",
-#if ROCM_VERSION >= 60300
-        "gfx1200", "gfx1201",
-#endif
-#if ROCM_VERSION >= 60500
-        "gfx950"
-#endif
+        "gfx90a", "gfx942"
     };
     for (auto index: c10::irange(detail::getCUDAHooks().deviceCount())) {
       if (!detail::getCUDAHooks().isGPUArch(index, archs)) {