initial gfx1036 and gfx1103 support

- initial support for gfx1036 and gfx1103 as a build target - updated also the gfx1010 configuration settings to be more similar in composable kernel and miopen fixes: #101 fixes: #103 Signed-off-by: Mika Laitio <[email protected]>
lamikr · Jul 14, 2024 · 4940337 · 4940337
1 parent 4ebd6df
commit 4940337
Show file tree

Hide file tree

Showing 32 changed files with 849 additions and 174 deletions.
diff --git a/patches/rocm-6.1.2/MIOpen/0001-Do-not-fail-on-install-for-missing-kdb.bz2-file.patch b/patches/rocm-6.1.2/MIOpen/0001-Do-not-fail-on-install-for-missing-kdb.bz2-file.patch
@@ -1,7 +1,7 @@
-From 915a3c074ff1c3694566fd74f2e83003a1253104 Mon Sep 17 00:00:00 2001
+From 7b51b6d06842de4851b6c20b61dda2ead78e68ff Mon Sep 17 00:00:00 2001
 From: Mika Laitio <[email protected]>
 Date: Wed, 8 May 2024 13:48:21 -0700
-Subject: [PATCH 1/4] Do not fail on install for missing kdb.bz2 file
+Subject: [PATCH 1/5] Do not fail on install for missing kdb.bz2 file
 
 Do not fail if kdb.bz2 file for some GPU does not exist because
 these kdb files are not mandatory. Their function is to speed up the
@@ -43,5 +43,5 @@ index 32d9a2e5b..d6c2db704 100644
  endif()
 
 -- 
-2.41.1
+2.45.2
 
diff --git a/patches/rocm-6.1.2/MIOpen/0002-fix-libroctx64.so-linking-error.patch b/patches/rocm-6.1.2/MIOpen/0002-fix-libroctx64.so-linking-error.patch
@@ -1,7 +1,7 @@
-From ecb981cd1b66749186404fa76c56237c758953b5 Mon Sep 17 00:00:00 2001
+From fdc0e268b2edbb0a57af8e8820b388816af38471 Mon Sep 17 00:00:00 2001
 From: Mika Laitio <[email protected]>
 Date: Fri, 3 May 2024 14:15:09 -0700
-Subject: [PATCH 2/4] fix libroctx64.so linking error
+Subject: [PATCH 2/5] fix libroctx64.so linking error
 
 search the library and if found link it from there
 instead of expecting it to be in the ld library path
@@ -31,5 +31,5 @@ index 0741a6023..ae4405eed 100644
 
  ############################################################
 -- 
-2.41.1
+2.45.2
 
diff --git a/patches/rocm-6.1.2/MIOpen/0003-MIOpen-gfx1010-and-gfx1035-support.patch b/patches/rocm-6.1.2/MIOpen/0003-MIOpen-gfx1010-and-gfx1035-support.patch
@@ -1,7 +1,7 @@
-From 6578a68e3226e97716aad12d445632358f2a463e Mon Sep 17 00:00:00 2001
+From 4b65dfcb0208bfe1eb64c474fb97e31e2b8bf60f Mon Sep 17 00:00:00 2001
 From: Mika Laitio <[email protected]>
 Date: Tue, 19 Dec 2023 15:13:46 -0800
-Subject: [PATCH 3/4] MIOpen gfx1010 and gfx1035 support
+Subject: [PATCH 3/5] MIOpen gfx1010 and gfx1035 support
 
 - todo: check gfx1010 specific parts
 
@@ -174,5 +174,5 @@ index 16ce78f04..2ec3eaf09 100644
                         "gfx1101",
                         "gfx1102"};
 -- 
-2.41.1
+2.45.2
 
diff --git a/patches/rocm-6.1.2/MIOpen/0004-improved-gfx1010-support.patch b/patches/rocm-6.1.2/MIOpen/0004-improved-gfx1010-support.patch
@@ -1,7 +1,7 @@
-From 5e7803271cbbe475da352ab188f09b345006d9c0 Mon Sep 17 00:00:00 2001
+From 3c0552223df66774e8f3613826a6094939763cce Mon Sep 17 00:00:00 2001
 From: Mika Laitio <[email protected]>
 Date: Mon, 8 Jul 2024 21:44:10 +0300
-Subject: [PATCH 4/4] improved gfx1010 support
+Subject: [PATCH 4/5] improved gfx1010 support
 
 - allows running pytorch gpu benchmark
   on gfx1010/amd rx 5700
@@ -135,5 +135,5 @@ index bf02d4d55..c3fa2bd3a 100644
          {"Rembrandt", "gfx1035"},
      };
 -- 
-2.41.1
+2.45.2
 
diff --git a/patches/rocm-6.1.2/MIOpen/0005-gfx1036-and-gfx1103-support.patch b/patches/rocm-6.1.2/MIOpen/0005-gfx1036-and-gfx1103-support.patch
@@ -0,0 +1,119 @@
+From 08071937d4c2c34f619ed5b49bd0ced4805875fa Mon Sep 17 00:00:00 2001
+From: Mika Laitio <[email protected]>
+Date: Sat, 13 Jul 2024 21:07:11 -0400
+Subject: [PATCH 5/5] gfx1036 and gfx1103 support
+
+Signed-off-by: Mika Laitio <[email protected]>
+---
+ .../composable_kernel/include/utility/config.hpp  | 15 ++++++++-------
+ src/include/miopen/solver/ck_utility_common.hpp   |  8 +++++++-
+ src/target_properties.cpp                         | 14 ++++++++++++--
+ 3 files changed, 27 insertions(+), 10 deletions(-)
+
+diff --git a/src/composable_kernel/composable_kernel/include/utility/config.hpp b/src/composable_kernel/composable_kernel/include/utility/config.hpp
+index 5957a79d8..6ca920b5e 100644
+--- a/src/composable_kernel/composable_kernel/include/utility/config.hpp
++++ b/src/composable_kernel/composable_kernel/include/utility/config.hpp
+@@ -16,8 +16,8 @@
+ #if !(defined(CK_AMD_GPU_GFX803) || defined(CK_AMD_GPU_GFX900) || defined(CK_AMD_GPU_GFX906) ||    \
+       defined(CK_AMD_GPU_GFX940) || defined(CK_AMD_GPU_GFX908) || defined(CK_AMD_GPU_GFX90A) ||    \
+       defined(CK_AMD_GPU_GFX941) || defined(CK_AMD_GPU_GFX942) || defined(CK_AMD_GPU_GFX1010) || defined(CK_AMD_GPU_GFX1030) ||   \
+-      defined(CK_AMD_GPU_GFX1031) || defined(CK_AMD_GPU_GFX1035) || defined(CK_AMD_GPU_GFX1100) || defined(CK_AMD_GPU_GFX1101) || \
+-      defined(CK_AMD_GPU_GFX1102))
++      defined(CK_AMD_GPU_GFX1031) || defined(CK_AMD_GPU_GFX1035) || defined(CK_AMD_GPU_GFX1036) || defined(CK_AMD_GPU_GFX1100) || defined(CK_AMD_GPU_GFX1101) || \
++      defined(CK_AMD_GPU_GFX1102) || defined(CK_AMD_GPU_GFX1103))
+ #error Need to define (only) one GPU target
+ #endif
+
+@@ -29,14 +29,15 @@
+ #define CK_MIN_BLOCK_PER_CU 2
+ #endif
+
+-// TODO: gfx1010 check CK_BUFFER_RESOURCE_3RD_DWORD
++// TODO: composable_kernel has differend CK_BUFFER_RESOURCE_3RD_DWORD for gfx110* devices
+ // buffer resourse
+ #if defined(CK_AMD_GPU_GFX803) || defined(CK_AMD_GPU_GFX900) || defined(CK_AMD_GPU_GFX906) || \
+     defined(CK_AMD_GPU_GFX941) || defined(CK_AMD_GPU_GFX942) || defined(CK_AMD_GPU_GFX940) || \
+-    defined(CK_AMD_GPU_GFX908) || defined(CK_AMD_GPU_GFX90A) || defined(CK_AMD_GPU_GFX1010)
++    defined(CK_AMD_GPU_GFX908) || defined(CK_AMD_GPU_GFX90A)
+ #define CK_BUFFER_RESOURCE_3RD_DWORD 0x00020000
+-#elif defined(CK_AMD_GPU_GFX1030) || defined(CK_AMD_GPU_GFX1031) || defined(CK_AMD_GPU_GFX1035) || defined(CK_AMD_GPU_GFX1100) || \
+-    defined(CK_AMD_GPU_GFX1101) || defined(CK_AMD_GPU_GFX1102)
++#elif defined(CK_AMD_GPU_GFX1010) || defined(CK_AMD_GPU_GFX1030) || defined(CK_AMD_GPU_GFX1031) || \
++    defined(CK_AMD_GPU_GFX1035) || defined(CK_AMD_GPU_GFX1036) || \
++    defined(CK_AMD_GPU_GFX1100) || defined(CK_AMD_GPU_GFX1101) || defined(CK_AMD_GPU_GFX1102) || defined(CK_AMD_GPU_GFX1103)
+ #define CK_BUFFER_RESOURCE_3RD_DWORD 0x31014000
+ #endif
+
+@@ -49,7 +50,7 @@
+ #elif defined(CK_AMD_GPU_GFX906) || defined(CK_AMD_GPU_GFX908) || defined(CK_AMD_GPU_GFX90a) ||  \
+     defined(CK_AMD_GPU_GFX941) || defined(CK_AMD_GPU_GFX942) || defined(CK_AMD_GPU_GFX940) ||    \
+     defined(CK_AMD_GPU_GFX1030) || defined(CK_AMD_GPU_GFX1031) || defined(CK_AMD_GPU_GFX1100) || \
+-    defined(CK_AMD_GPU_GFX1101) || defined(CK_AMD_GPU_GFX1102)
++    defined(CK_AMD_GPU_GFX1101) || defined(CK_AMD_GPU_GFX1102) || defined(CK_AMD_GPU_GFX1103)
+ #define CK_USE_AMD_V_FMAC_F32
+ #define CK_USE_AMD_V_DOT2_F32_F16
+ #define CK_USE_AMD_V_DOT4_I32_I8
+diff --git a/src/include/miopen/solver/ck_utility_common.hpp b/src/include/miopen/solver/ck_utility_common.hpp
+index aea036066..ea5629871 100644
+--- a/src/include/miopen/solver/ck_utility_common.hpp
++++ b/src/include/miopen/solver/ck_utility_common.hpp
+@@ -61,9 +61,11 @@ static inline bool is_ck_supported_hardware(const Handle& handle)
+            StartsWith(handle.GetDeviceName(), "gfx1030") ||
+            StartsWith(handle.GetDeviceName(), "gfx1031") ||
+            StartsWith(handle.GetDeviceName(), "gfx1035") ||
++           StartsWith(handle.GetDeviceName(), "gfx1036") ||
+            StartsWith(handle.GetDeviceName(), "gfx1100") ||
+            StartsWith(handle.GetDeviceName(), "gfx1101") ||
+-           StartsWith(handle.GetDeviceName(), "gfx1102");
++           StartsWith(handle.GetDeviceName(), "gfx1102") ||
++           StartsWith(handle.GetDeviceName(), "gfx1103");
+ }
+
+ // MI100 : gfx908
+@@ -121,12 +123,16 @@ static inline auto get_ck_common_compiler_flag(const Handle& handle)
+         compiler_flag << " -DCK_AMD_GPU_GFX1031";
+     else if(StartsWith(device_name, "gfx1035"))
+         compiler_flag << " -DCK_AMD_GPU_GFX1035";
++    else if(StartsWith(device_name, "gfx1036"))
++        compiler_flag << " -DCK_AMD_GPU_GFX1036";
+     else if(StartsWith(device_name, "gfx1100"))
+         compiler_flag << " -DCK_AMD_GPU_GFX1100";
+     else if(StartsWith(device_name, "gfx1101"))
+         compiler_flag << " -DCK_AMD_GPU_GFX1101";
+     else if(StartsWith(device_name, "gfx1102"))
+         compiler_flag << " -DCK_AMD_GPU_GFX1102";
++    else if(StartsWith(device_name, "gfx1103"))
++        compiler_flag << " -DCK_AMD_GPU_GFX1103";
+     // NOLINTEND(*-braces-around-statements)
+
+     // buffer atomic-fadd
+diff --git a/src/target_properties.cpp b/src/target_properties.cpp
+index c3fa2bd3a..de979aae9 100644
+--- a/src/target_properties.cpp
++++ b/src/target_properties.cpp
+@@ -52,9 +52,19 @@ static std::string GetDeviceNameFromMap(const std::string& in)
+         {"gfx804", "gfx803"},
+         {"Vega10", "gfx900"},
+         {"gfx901", "gfx900"},
+-        {"Navi10", "gfx1010"},
++        {"navi10", "gfx1010"},
++        {"navi12", "gfx1011"},
++        {"navi14", "gfx1012"},
+         {"10.3.0 Sienna_Cichlid 18", "gfx1030"},
+-        {"Rembrandt", "gfx1035"},
++        {"navi22", "gfx1031"},
++        {"navi23", "gfx1032"},
++        {"navi24", "gfx1034"},
++        {"rembrandt", "gfx1035"},
++        {"rembrandt1036", "gfx1036"},
++        {"navi31", "gfx1100"},
++        {"navi32", "gfx1101"},
++        {"navi33", "gfx1102"},
++        {"phoenix", "gfx1103"},
+     };
+
+     const auto& dev_str = miopen::GetStringEnv(ENV(MIOPEN_DEBUG_ENFORCE_DEVICE));
+-- 
+2.45.2
+
diff --git a/patches/rocm-6.1.2/Tensile/0001-Tensile-fix-fallback-arch-build.patch b/patches/rocm-6.1.2/Tensile/0001-Tensile-fix-fallback-arch-build.patch
@@ -1,7 +1,7 @@
-From dd39290f90df98f928452720e622c1497b4fd7da Mon Sep 17 00:00:00 2001
+From 3e2d2f891001a8d2a8f74a46884cbec84a2fd8c8 Mon Sep 17 00:00:00 2001
 From: Mika Laitio <[email protected]>
 Date: Fri, 10 May 2024 20:34:13 -0700
-Subject: [PATCH 1/3] Tensile fix fallback arch build
+Subject: [PATCH 1/4] Tensile fix fallback arch build
 
 fixes build error which happens if only the rx 5700 is enabled
 (only GPU_BUILD_AMD_NAVI10_GFX1010=1 enabled in rocm_sdk_builder envsetup.sh)
@@ -61,5 +61,5 @@ index ca3ef322..9e37b4b0 100644
    for arch in archs:
      if arch in architectureMap:
 -- 
-2.41.0
+2.45.2
 
diff --git a/patches/rocm-6.1.2/Tensile/0002-Tensile-add-gfx1035-support.patch b/patches/rocm-6.1.2/Tensile/0002-Tensile-add-gfx1035-support.patch
@@ -1,7 +1,7 @@
-From 73281f46189f7012334d3b1a7e52baffade5295f Mon Sep 17 00:00:00 2001
+From cbec649b52abb1e45d72a3755fa01a77cf9784e7 Mon Sep 17 00:00:00 2001
 From: Mika Laitio <[email protected]>
 Date: Fri, 3 May 2024 13:13:02 -0700
-Subject: [PATCH 2/3] Tensile, add gfx1035 support
+Subject: [PATCH 2/4] Tensile, add gfx1035 support
 
 Signed-off-by: Mika Laitio <[email protected]>
 ---
@@ -84,5 +84,5 @@ index 6ececf1c..a89b7c39 100644
    else:
      printWarning("Assembler not present, asm caps loaded from cache are unverified")
 -- 
-2.41.0
+2.45.2
 
diff --git a/patches/rocm-6.1.2/Tensile/0003-llvm-path-changes.patch b/patches/rocm-6.1.2/Tensile/0003-llvm-path-changes.patch
@@ -1,7 +1,7 @@
-From 3ffffba383528d997372861fb5d940d09ebe2996 Mon Sep 17 00:00:00 2001
+From 80776357e8f44019675a224474e314bbb551bc8a Mon Sep 17 00:00:00 2001
 From: Mika Laitio <[email protected]>
 Date: Wed, 15 May 2024 21:09:56 -0700
-Subject: [PATCH 3/3] llvm path changes
+Subject: [PATCH 3/4] llvm path changes
 
 Signed-off-by: Mika Laitio <[email protected]>
 ---
@@ -18,17 +18,18 @@ index a89b7c39..9b24aee7 100644
      else:
 -      globalParameters["AssemblerPath"] = locateExe(os.path.join(globalParameters["ROCmPath"], "llvm/bin"), "clang++")
 +      globalParameters["AssemblerPath"] = locateExe(os.path.join(globalParameters["ROCmPath"], "bin"), "clang++")
-
+ 
    globalParameters["ROCmSMIPath"] = locateExe(globalParameters["ROCmBinPath"], "rocm-smi")
-
+ 
 @@ -2277,7 +2277,7 @@ def assignGlobalParameters( config ):
      if os.name == "nt":
        globalParameters["ClangOffloadBundlerPath"] = locateExe(globalParameters["ROCmBinPath"], "clang-offload-bundler.exe")
      else:
 -      globalParameters["ClangOffloadBundlerPath"] = locateExe(os.path.join(globalParameters["ROCmPath"], "llvm/bin"), "clang-offload-bundler")
 +      globalParameters["ClangOffloadBundlerPath"] = locateExe(os.path.join(globalParameters["ROCmPath"], "bin"), "clang-offload-bundler")
-
+ 
    if "ROCmAgentEnumeratorPath" in config:
      globalParameters["ROCmAgentEnumeratorPath"] = config["ROCmAgentEnumeratorPath"]
---
-2.41.0
+-- 
+2.45.2
+