[ROCm] enable gfx110x for hipblaslt (pytorch#137317)

Fixes pytorch#136347.

Pull Request resolved: pytorch#137317
Approved by: https://github.com/Skylion007, https://github.com/jithunnair-amd

Co-authored-by: Nichols A. Romero <[email protected]>

Author: 2 people

aten/src/ATen/Context.cpp

@@ -290,7 +290,12 @@ at::BlasBackend Context::blasPreferredBackend() {
 #ifdef USE_ROCM
   if (blas_preferred_backend == at::BlasBackend::Cublaslt) {
     static const bool hipblaslt_unsupported = []() {
-      static const std::vector<std::string> archs = {"gfx90a", "gfx940", "gfx941", "gfx942"};
+      static const std::vector<std::string> archs = {
+          "gfx90a", "gfx940", "gfx941", "gfx942",
+#if ROCM_VERSION >= 60300
+          "gfx1100", "gfx1101"
+#endif
+      };
       for (auto index: c10::irange(getNumGPUs())) {
         if (!detail::getCUDAHooks().isGPUArch(index, archs)) {
           TORCH_WARN_ONCE(

aten/src/ATen/native/cuda/Blas.cpp

@@ -201,7 +201,12 @@ static bool getDisableAddmmCudaLt() {
 static bool isSupportedHipLtROCmArch(int index) {
     hipDeviceProp_t* prop = at::cuda::getDeviceProperties(index);
     std::string device_arch = prop->gcnArchName;
-    static const std::vector<std::string> archs = {"gfx90a", "gfx940", "gfx941", "gfx942"};
+    static const std::vector<std::string> archs = {
+        "gfx90a", "gfx940", "gfx941", "gfx942",
+#if ROCM_VERSION >= 60300
+        "gfx1100", "gfx1101"
+#endif
+    };
     for (std::string arch : archs) {
         size_t substring = device_arch.find(arch);
         if (substring != std::string::npos) {