Skip to content

Commit 1e191b4

Browse files
lamikrlamikr
committed
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 <lamikr@gmail.com>
1 parent c62cca1 commit 1e191b4

File tree

47 files changed

+1076
-257
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

47 files changed

+1076
-257
lines changed

‎binfo/user_config.sh‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -3,7 +3,7 @@
33
# Function to select ROCM SDK build target GPUs
44
func_build_cfg_user() {
55
local message="Select ROCM SDK build target GPUs. Space to select, Enter to finish save, ESC to cancel."
6-
local options="gfx906|gfx90a|gfx940|gfx1010|gfx1011|gfx1012|gfx1030|gfx1031|gfx1035|gfx1100|gfx1101|gfx1102|gfx1150|gfx1151"
6+
local options="gfx906|gfx90a|gfx940|gfx1010|gfx1011|gfx1012|gfx1030|gfx1031|gfx1035|gfx1036|gfx1100|gfx1101|gfx1102|gfx1103|gfx1150|gfx1151"
77
local script_path="./build/checkbox.sh"
88

99
# Check if the script exists and is executable
@@ -14,4 +14,4 @@ func_build_cfg_user() {
1414

1515
# Execute the checkbox script with the specified parameters
1616
"$script_path" --message="$message" --options="$options" --multiple
17-
}
17+
}

‎patches/rocm-6.1.2/MIOpen/0001-Do-not-fail-on-install-for-missing-kdb.bz2-file.patch‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
1-
From 915a3c074ff1c3694566fd74f2e83003a1253104 Mon Sep 17 00:00:00 2001
1+
From 7b51b6d06842de4851b6c20b61dda2ead78e68ff Mon Sep 17 00:00:00 2001
22
From: Mika Laitio <lamikr@pilppa.org>
33
Date: Wed, 8 May 2024 13:48:21 -0700
4-
Subject: [PATCH 1/4] Do not fail on install for missing kdb.bz2 file
4+
Subject: [PATCH 1/5] Do not fail on install for missing kdb.bz2 file
55

66
Do not fail if kdb.bz2 file for some GPU does not exist because
77
these kdb files are not mandatory. Their function is to speed up the
@@ -43,5 +43,5 @@ index 32d9a2e5b..d6c2db704 100644
4343
endif()
4444

4545
--
46-
2.41.1
46+
2.45.2
4747

‎patches/rocm-6.1.2/MIOpen/0002-fix-libroctx64.so-linking-error.patch‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
1-
From ecb981cd1b66749186404fa76c56237c758953b5 Mon Sep 17 00:00:00 2001
1+
From fdc0e268b2edbb0a57af8e8820b388816af38471 Mon Sep 17 00:00:00 2001
22
From: Mika Laitio <lamikr@pilppa.org>
33
Date: Fri, 3 May 2024 14:15:09 -0700
4-
Subject: [PATCH 2/4] fix libroctx64.so linking error
4+
Subject: [PATCH 2/5] fix libroctx64.so linking error
55

66
search the library and if found link it from there
77
instead of expecting it to be in the ld library path
@@ -31,5 +31,5 @@ index 0741a6023..ae4405eed 100644
3131

3232
############################################################
3333
--
34-
2.41.1
34+
2.45.2
3535

‎patches/rocm-6.1.2/MIOpen/0003-MIOpen-gfx1010-and-gfx1035-support.patch‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
1-
From 6578a68e3226e97716aad12d445632358f2a463e Mon Sep 17 00:00:00 2001
1+
From 4b65dfcb0208bfe1eb64c474fb97e31e2b8bf60f Mon Sep 17 00:00:00 2001
22
From: Mika Laitio <lamikr@gmail.com>
33
Date: Tue, 19 Dec 2023 15:13:46 -0800
4-
Subject: [PATCH 3/4] MIOpen gfx1010 and gfx1035 support
4+
Subject: [PATCH 3/5] MIOpen gfx1010 and gfx1035 support
55

66
- todo: check gfx1010 specific parts
77

@@ -174,5 +174,5 @@ index 16ce78f04..2ec3eaf09 100644
174174
"gfx1101",
175175
"gfx1102"};
176176
--
177-
2.41.1
177+
2.45.2
178178

‎patches/rocm-6.1.2/MIOpen/0004-improved-gfx1010-support.patch‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
1-
From 5e7803271cbbe475da352ab188f09b345006d9c0 Mon Sep 17 00:00:00 2001
1+
From 3c0552223df66774e8f3613826a6094939763cce Mon Sep 17 00:00:00 2001
22
From: Mika Laitio <lamikr@pilppa.org>
33
Date: Mon, 8 Jul 2024 21:44:10 +0300
4-
Subject: [PATCH 4/4] improved gfx1010 support
4+
Subject: [PATCH 4/5] improved gfx1010 support
55

66
- allows running pytorch gpu benchmark
77
on gfx1010/amd rx 5700
@@ -135,5 +135,5 @@ index bf02d4d55..c3fa2bd3a 100644
135135
{"Rembrandt", "gfx1035"},
136136
};
137137
--
138-
2.41.1
138+
2.45.2
139139

‎patches/rocm-6.1.2/MIOpen/0005-gfx1036-and-gfx1103-support.patch‎

Lines changed: 119 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,119 @@
1+
From 08071937d4c2c34f619ed5b49bd0ced4805875fa Mon Sep 17 00:00:00 2001
2+
From: Mika Laitio <lamikr@gmail.com>
3+
Date: Sat, 13 Jul 2024 21:07:11 -0400
4+
Subject: [PATCH 5/5] gfx1036 and gfx1103 support
5+
6+
Signed-off-by: Mika Laitio <lamikr@gmail.com>
7+
---
8+
.../composable_kernel/include/utility/config.hpp | 15 ++++++++-------
9+
src/include/miopen/solver/ck_utility_common.hpp | 8 +++++++-
10+
src/target_properties.cpp | 14 ++++++++++++--
11+
3 files changed, 27 insertions(+), 10 deletions(-)
12+
13+
diff --git a/src/composable_kernel/composable_kernel/include/utility/config.hpp b/src/composable_kernel/composable_kernel/include/utility/config.hpp
14+
index 5957a79d8..6ca920b5e 100644
15+
--- a/src/composable_kernel/composable_kernel/include/utility/config.hpp
16+
+++ b/src/composable_kernel/composable_kernel/include/utility/config.hpp
17+
@@ -16,8 +16,8 @@
18+
#if !(defined(CK_AMD_GPU_GFX803) || defined(CK_AMD_GPU_GFX900) || defined(CK_AMD_GPU_GFX906) || \
19+
defined(CK_AMD_GPU_GFX940) || defined(CK_AMD_GPU_GFX908) || defined(CK_AMD_GPU_GFX90A) || \
20+
defined(CK_AMD_GPU_GFX941) || defined(CK_AMD_GPU_GFX942) || defined(CK_AMD_GPU_GFX1010) || defined(CK_AMD_GPU_GFX1030) || \
21+
- defined(CK_AMD_GPU_GFX1031) || defined(CK_AMD_GPU_GFX1035) || defined(CK_AMD_GPU_GFX1100) || defined(CK_AMD_GPU_GFX1101) || \
22+
- defined(CK_AMD_GPU_GFX1102))
23+
+ 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) || \
24+
+ defined(CK_AMD_GPU_GFX1102) || defined(CK_AMD_GPU_GFX1103))
25+
#error Need to define (only) one GPU target
26+
#endif
27+
28+
@@ -29,14 +29,15 @@
29+
#define CK_MIN_BLOCK_PER_CU 2
30+
#endif
31+
32+
-// TODO: gfx1010 check CK_BUFFER_RESOURCE_3RD_DWORD
33+
+// TODO: composable_kernel has differend CK_BUFFER_RESOURCE_3RD_DWORD for gfx110* devices
34+
// buffer resourse
35+
#if defined(CK_AMD_GPU_GFX803) || defined(CK_AMD_GPU_GFX900) || defined(CK_AMD_GPU_GFX906) || \
36+
defined(CK_AMD_GPU_GFX941) || defined(CK_AMD_GPU_GFX942) || defined(CK_AMD_GPU_GFX940) || \
37+
- defined(CK_AMD_GPU_GFX908) || defined(CK_AMD_GPU_GFX90A) || defined(CK_AMD_GPU_GFX1010)
38+
+ defined(CK_AMD_GPU_GFX908) || defined(CK_AMD_GPU_GFX90A)
39+
#define CK_BUFFER_RESOURCE_3RD_DWORD 0x00020000
40+
-#elif defined(CK_AMD_GPU_GFX1030) || defined(CK_AMD_GPU_GFX1031) || defined(CK_AMD_GPU_GFX1035) || defined(CK_AMD_GPU_GFX1100) || \
41+
- defined(CK_AMD_GPU_GFX1101) || defined(CK_AMD_GPU_GFX1102)
42+
+#elif defined(CK_AMD_GPU_GFX1010) || defined(CK_AMD_GPU_GFX1030) || defined(CK_AMD_GPU_GFX1031) || \
43+
+ defined(CK_AMD_GPU_GFX1035) || defined(CK_AMD_GPU_GFX1036) || \
44+
+ defined(CK_AMD_GPU_GFX1100) || defined(CK_AMD_GPU_GFX1101) || defined(CK_AMD_GPU_GFX1102) || defined(CK_AMD_GPU_GFX1103)
45+
#define CK_BUFFER_RESOURCE_3RD_DWORD 0x31014000
46+
#endif
47+
48+
@@ -49,7 +50,7 @@
49+
#elif defined(CK_AMD_GPU_GFX906) || defined(CK_AMD_GPU_GFX908) || defined(CK_AMD_GPU_GFX90a) || \
50+
defined(CK_AMD_GPU_GFX941) || defined(CK_AMD_GPU_GFX942) || defined(CK_AMD_GPU_GFX940) || \
51+
defined(CK_AMD_GPU_GFX1030) || defined(CK_AMD_GPU_GFX1031) || defined(CK_AMD_GPU_GFX1100) || \
52+
- defined(CK_AMD_GPU_GFX1101) || defined(CK_AMD_GPU_GFX1102)
53+
+ defined(CK_AMD_GPU_GFX1101) || defined(CK_AMD_GPU_GFX1102) || defined(CK_AMD_GPU_GFX1103)
54+
#define CK_USE_AMD_V_FMAC_F32
55+
#define CK_USE_AMD_V_DOT2_F32_F16
56+
#define CK_USE_AMD_V_DOT4_I32_I8
57+
diff --git a/src/include/miopen/solver/ck_utility_common.hpp b/src/include/miopen/solver/ck_utility_common.hpp
58+
index aea036066..ea5629871 100644
59+
--- a/src/include/miopen/solver/ck_utility_common.hpp
60+
+++ b/src/include/miopen/solver/ck_utility_common.hpp
61+
@@ -61,9 +61,11 @@ static inline bool is_ck_supported_hardware(const Handle& handle)
62+
StartsWith(handle.GetDeviceName(), "gfx1030") ||
63+
StartsWith(handle.GetDeviceName(), "gfx1031") ||
64+
StartsWith(handle.GetDeviceName(), "gfx1035") ||
65+
+ StartsWith(handle.GetDeviceName(), "gfx1036") ||
66+
StartsWith(handle.GetDeviceName(), "gfx1100") ||
67+
StartsWith(handle.GetDeviceName(), "gfx1101") ||
68+
- StartsWith(handle.GetDeviceName(), "gfx1102");
69+
+ StartsWith(handle.GetDeviceName(), "gfx1102") ||
70+
+ StartsWith(handle.GetDeviceName(), "gfx1103");
71+
}
72+
73+
// MI100 : gfx908
74+
@@ -121,12 +123,16 @@ static inline auto get_ck_common_compiler_flag(const Handle& handle)
75+
compiler_flag << " -DCK_AMD_GPU_GFX1031";
76+
else if(StartsWith(device_name, "gfx1035"))
77+
compiler_flag << " -DCK_AMD_GPU_GFX1035";
78+
+ else if(StartsWith(device_name, "gfx1036"))
79+
+ compiler_flag << " -DCK_AMD_GPU_GFX1036";
80+
else if(StartsWith(device_name, "gfx1100"))
81+
compiler_flag << " -DCK_AMD_GPU_GFX1100";
82+
else if(StartsWith(device_name, "gfx1101"))
83+
compiler_flag << " -DCK_AMD_GPU_GFX1101";
84+
else if(StartsWith(device_name, "gfx1102"))
85+
compiler_flag << " -DCK_AMD_GPU_GFX1102";
86+
+ else if(StartsWith(device_name, "gfx1103"))
87+
+ compiler_flag << " -DCK_AMD_GPU_GFX1103";
88+
// NOLINTEND(*-braces-around-statements)
89+
90+
// buffer atomic-fadd
91+
diff --git a/src/target_properties.cpp b/src/target_properties.cpp
92+
index c3fa2bd3a..de979aae9 100644
93+
--- a/src/target_properties.cpp
94+
+++ b/src/target_properties.cpp
95+
@@ -52,9 +52,19 @@ static std::string GetDeviceNameFromMap(const std::string& in)
96+
{"gfx804", "gfx803"},
97+
{"Vega10", "gfx900"},
98+
{"gfx901", "gfx900"},
99+
- {"Navi10", "gfx1010"},
100+
+ {"navi10", "gfx1010"},
101+
+ {"navi12", "gfx1011"},
102+
+ {"navi14", "gfx1012"},
103+
{"10.3.0 Sienna_Cichlid 18", "gfx1030"},
104+
- {"Rembrandt", "gfx1035"},
105+
+ {"navi22", "gfx1031"},
106+
+ {"navi23", "gfx1032"},
107+
+ {"navi24", "gfx1034"},
108+
+ {"rembrandt", "gfx1035"},
109+
+ {"raphael", "gfx1036"},
110+
+ {"navi31", "gfx1100"},
111+
+ {"navi32", "gfx1101"},
112+
+ {"navi33", "gfx1102"},
113+
+ {"phoenix", "gfx1103"},
114+
};
115+
116+
const auto& dev_str = miopen::GetStringEnv(ENV(MIOPEN_DEBUG_ENFORCE_DEVICE));
117+
--
118+
2.45.2
119+

‎patches/rocm-6.1.2/Tensile/0001-Tensile-fix-fallback-arch-build.patch‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
1-
From dd39290f90df98f928452720e622c1497b4fd7da Mon Sep 17 00:00:00 2001
1+
From 3e2d2f891001a8d2a8f74a46884cbec84a2fd8c8 Mon Sep 17 00:00:00 2001
22
From: Mika Laitio <lamikr@pilppa.org>
33
Date: Fri, 10 May 2024 20:34:13 -0700
4-
Subject: [PATCH 1/3] Tensile fix fallback arch build
4+
Subject: [PATCH 1/4] Tensile fix fallback arch build
55

66
fixes build error which happens if only the rx 5700 is enabled
77
(only GPU_BUILD_AMD_NAVI10_GFX1010=1 enabled in rocm_sdk_builder envsetup.sh)
@@ -61,5 +61,5 @@ index ca3ef322..9e37b4b0 100644
6161
for arch in archs:
6262
if arch in architectureMap:
6363
--
64-
2.41.0
64+
2.45.2
6565

‎patches/rocm-6.1.2/Tensile/0002-Tensile-add-gfx1035-support.patch‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
1-
From 73281f46189f7012334d3b1a7e52baffade5295f Mon Sep 17 00:00:00 2001
1+
From cbec649b52abb1e45d72a3755fa01a77cf9784e7 Mon Sep 17 00:00:00 2001
22
From: Mika Laitio <lamikr@pilppa.org>
33
Date: Fri, 3 May 2024 13:13:02 -0700
4-
Subject: [PATCH 2/3] Tensile, add gfx1035 support
4+
Subject: [PATCH 2/4] Tensile, add gfx1035 support
55

66
Signed-off-by: Mika Laitio <lamikr@pilppa.org>
77
---
@@ -84,5 +84,5 @@ index 6ececf1c..a89b7c39 100644
8484
else:
8585
printWarning("Assembler not present, asm caps loaded from cache are unverified")
8686
--
87-
2.41.0
87+
2.45.2
8888

‎patches/rocm-6.1.2/Tensile/0003-llvm-path-changes.patch‎

Lines changed: 8 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
1-
From 3ffffba383528d997372861fb5d940d09ebe2996 Mon Sep 17 00:00:00 2001
1+
From 80776357e8f44019675a224474e314bbb551bc8a Mon Sep 17 00:00:00 2001
22
From: Mika Laitio <lamikr@gmail.com>
33
Date: Wed, 15 May 2024 21:09:56 -0700
4-
Subject: [PATCH 3/3] llvm path changes
4+
Subject: [PATCH 3/4] llvm path changes
55

66
Signed-off-by: Mika Laitio <lamikr@gmail.com>
77
---
@@ -18,17 +18,18 @@ index a89b7c39..9b24aee7 100644
1818
else:
1919
- globalParameters["AssemblerPath"] = locateExe(os.path.join(globalParameters["ROCmPath"], "llvm/bin"), "clang++")
2020
+ globalParameters["AssemblerPath"] = locateExe(os.path.join(globalParameters["ROCmPath"], "bin"), "clang++")
21-
21+
2222
globalParameters["ROCmSMIPath"] = locateExe(globalParameters["ROCmBinPath"], "rocm-smi")
23-
23+
2424
@@ -2277,7 +2277,7 @@ def assignGlobalParameters( config ):
2525
if os.name == "nt":
2626
globalParameters["ClangOffloadBundlerPath"] = locateExe(globalParameters["ROCmBinPath"], "clang-offload-bundler.exe")
2727
else:
2828
- globalParameters["ClangOffloadBundlerPath"] = locateExe(os.path.join(globalParameters["ROCmPath"], "llvm/bin"), "clang-offload-bundler")
2929
+ globalParameters["ClangOffloadBundlerPath"] = locateExe(os.path.join(globalParameters["ROCmPath"], "bin"), "clang-offload-bundler")
30-
30+
3131
if "ROCmAgentEnumeratorPath" in config:
3232
globalParameters["ROCmAgentEnumeratorPath"] = config["ROCmAgentEnumeratorPath"]
33-
--
34-
2.41.0
33+
--
34+
2.45.2
35+

0 commit comments

Comments
 (0)