summaryrefslogtreecommitdiff
path: root/sci-libs/miopen/files/miopen-10.1.0-llvm-23-asm-fake16.patch
blob: 3238140a8fcac7c7d5a0f4a9474660cc5b88ad12 (plain)
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
Fix assembling MIOpen asm kernels for gfx11/gfx12 with LLVM 23

LLVM 23 enables real-true16 on gfx11/gfx12 and rejects the fake16 syntax
(e.g. v_cvt_f16_f32 v1, s32) used by the Winograd asm kernels, which comgr
assembles at runtime.
--- a/src/comgr.cpp
+++ b/src/comgr.cpp
@@ -514,6 +514,12 @@
             optAsm.emplace_back("-mno-xnack");
 #endif
         compiler::lc::gcnasm::RemoveOptionsUnwanted(optAsm);
+        // The asm kernels use fake16 syntax, which LLVM 23 rejects in real-true16 mode.
+        if(StartsWith(target.Name(), "gfx11") || StartsWith(target.Name(), "gfx12"))
+        {
+            for(const auto* opt : {"-Xclangas", "-target-feature", "-Xclangas", "-real-true16"})
+                optAsm.emplace_back(opt);
+        }
 
         action.SetOptionList(optAsm);