[llvm-branch-commits] [clang] [llvm] AMDGPU: Use module flags to control xnack and sramecc (PR #204595)
Matt Arsenault via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Thu Jun 18 06:38:21 PDT 2026
https://github.com/arsenm created https://github.com/llvm/llvm-project/pull/204595
This ensures these ABI details are encoded in the IR module
rather than depending on external state from command-line flags.
Previously, these were encoded as function-level subtarget features.
The code object output was a single target ID directive implied
by the global subtarget. The backend would previously check if a
function's subtarget feature mismatched the global subtarget. This
is avoided by making xnack and sramecc module-level properties from
the start. This also provides proper linker compatibility
enforcement, moving the error point earlier.
The old encoding was also an abuse of the subtarget feature system.
Subtarget features are a bitvector, and later features in the string
can override earlier ones. The old handling added a special case
where explicit settings were preserved: ordinarily +feature,-feature
should result in the feature being disabled, but +xnack,-xnack would
preserve the explicit "-xnack" state, which differs from the absence
of any xnack setting.
The new flags are encoded as 0/1, with the "any" case represented
as the absence of the flag. I considered an explicit tri-state unknown
value, but decided against it.
This also removes warnings when using these module flags on targets
that do not support the corresponding feature. Previously, messages
were written directly to stderr instead of using proper diagnostics.
Avoiding the warning reduces burden on frontends to check which targets
require the flags.
For migration purposes, the subtarget features still exist. Currently,
they are still respected in the various binary tools, pending
disassembler changes to determine target ID modifiers from e_flags.
CodeGen requires using the module flags. An error will be raised when
attempting to use the old global subtarget features. These should be
removed after a migration period for frontends to update. Functionality
wise, bitcode autoupgrade should work. Old bitcode will not have the
flags, resulting in a different target ID in the output binary than
expected, but it should run correctly.
New cl::opts exist only because it was inconvenient to update all
tests using multiple xnack modes. Users should never use these.
Co-Authored-By: Claude Opus 4.6 <noreply at anthropic.com>
>From 8e181f46a7d69e2a2aac10d570e904b47fba82de Mon Sep 17 00:00:00 2001
From: Matt Arsenault <Matthew.Arsenault at amd.com>
Date: Sun, 14 Jun 2026 09:55:39 +0200
Subject: [PATCH] AMDGPU: Use module flags to control xnack and sramecc
This ensures these ABI details are encoded in the IR module
rather than depending on external state from command-line flags.
Previously, these were encoded as function-level subtarget features.
The code object output was a single target ID directive implied
by the global subtarget. The backend would previously check if a
function's subtarget feature mismatched the global subtarget. This
is avoided by making xnack and sramecc module-level properties from
the start. This also provides proper linker compatibility
enforcement, moving the error point earlier.
The old encoding was also an abuse of the subtarget feature system.
Subtarget features are a bitvector, and later features in the string
can override earlier ones. The old handling added a special case
where explicit settings were preserved: ordinarily +feature,-feature
should result in the feature being disabled, but +xnack,-xnack would
preserve the explicit "-xnack" state, which differs from the absence
of any xnack setting.
The new flags are encoded as 0/1, with the "any" case represented
as the absence of the flag. I considered an explicit tri-state unknown
value, but decided against it.
This also removes warnings when using these module flags on targets
that do not support the corresponding feature. Previously, messages
were written directly to stderr instead of using proper diagnostics.
Avoiding the warning reduces burden on frontends to check which targets
require the flags.
For migration purposes, the subtarget features still exist. Currently,
they are still respected in the various binary tools, pending
disassembler changes to determine target ID modifiers from e_flags.
CodeGen requires using the module flags. An error will be raised when
attempting to use the old global subtarget features. These should be
removed after a migration period for frontends to update. Functionality
wise, bitcode autoupgrade should work. Old bitcode will not have the
flags, resulting in a different target ID in the output binary than
expected, but it should run correctly.
New cl::opts exist only because it was inconvenient to update all
tests using multiple xnack modes. Users should never use these.
Co-Authored-By: Claude Opus 4.6 <noreply at anthropic.com>
---
clang/include/clang/Basic/TargetOptions.h | 16 +++
clang/include/clang/Options/Options.td | 16 ++-
clang/lib/Basic/Targets/AMDGPU.cpp | 10 ++
clang/lib/CodeGen/CodeGenModule.cpp | 21 ++++
clang/lib/Driver/ToolChains/Clang.cpp | 13 +++
clang/lib/Driver/ToolChains/Clang.h | 2 +
clang/lib/Frontend/CompilerInvocation.cpp | 27 +++++
clang/test/CIR/CodeGenHIP/target-features.hip | 17 +--
.../CodeGenCXX/dynamic-cast-address-space.cpp | 4 +-
.../amdgpu-module-flag-xnack-sramecc.cl | 22 ++++
.../CodeGenOpenCL/amdgpu-xnack-any-only.cl | 27 +++++
clang/test/Driver/amdgpu-features.c | 8 +-
clang/test/Driver/amdgpu-openmp-toolchain.c | 2 +-
clang/test/Driver/amdgpu-toolchain.c | 4 +-
.../test/Driver/amdgpu-xnack-sramecc-flags.c | 69 ++++++------
clang/test/Driver/hip-sanitize-options.hip | 34 +++---
clang/test/Driver/hip-target-id.hip | 14 +--
clang/test/Driver/hip-toolchain-features.hip | 16 +--
clang/test/Driver/target-id.cl | 10 +-
llvm/docs/AMDGPUUsage.rst | 40 ++++++-
llvm/docs/ReleaseNotes.md | 2 +
llvm/lib/IR/Verifier.cpp | 19 ++++
llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp | 82 ++++++--------
.../lib/Target/AMDGPU/AMDGPUTargetMachine.cpp | 46 +++++++-
llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.h | 5 +
.../AMDGPU/AsmParser/AMDGPUAsmParser.cpp | 2 +-
llvm/lib/Target/AMDGPU/GCNSubtarget.cpp | 27 +++--
llvm/lib/Target/AMDGPU/GCNSubtarget.h | 6 +-
.../MCTargetDesc/AMDGPUTargetStreamer.cpp | 8 ++
.../MCTargetDesc/AMDGPUTargetStreamer.h | 6 +-
.../Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp | 51 ++++++++-
llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h | 13 +++
.../GlobalISel/extractelement-stack-lower.ll | 2 +-
llvm/test/CodeGen/AMDGPU/amdpal-callable.ll | 6 +-
.../AMDGPU/break-smem-soft-clauses.mir | 4 +-
.../callee-special-input-vgprs-packed.ll | 2 +-
.../AMDGPU/cluster-flat-loads-postra.mir | 4 +-
llvm/test/CodeGen/AMDGPU/cluster_stores.ll | 2 +-
.../CodeGen/AMDGPU/directive-amdgcn-target.ll | 104 +++++++++---------
.../AMDGPU/elf-header-flags-sramecc.ll | 16 +--
.../CodeGen/AMDGPU/elf-header-flags-xnack.ll | 4 +-
llvm/test/CodeGen/AMDGPU/flat-saddr-load.ll | 4 +-
llvm/test/CodeGen/AMDGPU/flat-scratch-reg.ll | 22 ++--
.../CodeGen/AMDGPU/gfx902-without-xnack.ll | 5 +-
.../AMDGPU/greedy-reverse-local-assignment.ll | 4 +-
.../CodeGen/AMDGPU/hazard-hidden-bundle.mir | 6 +-
llvm/test/CodeGen/AMDGPU/hazard-in-bundle.mir | 6 +-
.../AMDGPU/hsa-metadata-kernel-code-props.ll | 6 +-
...tadata-resource-usage-function-ordering.ll | 6 +-
llvm/test/CodeGen/AMDGPU/hsa-note-no-func.ll | 8 +-
llvm/test/CodeGen/AMDGPU/immv216.ll | 6 +-
.../AMDGPU/limit-soft-clause-reg-pressure.mir | 4 +-
.../AMDGPU/materialize-frame-index-sgpr.ll | 6 +-
.../AMDGPU/mattr-xnack-sramecc-legacy.ll | 23 ++++
.../CodeGen/AMDGPU/module-flag-sramecc.ll | 66 +++++++++++
.../module-flag-xnack-sramecc-combined.ll | 15 +++
llvm/test/CodeGen/AMDGPU/module-flag-xnack.ll | 75 +++++++++++++
llvm/test/CodeGen/AMDGPU/nsa-reassign.ll | 2 +-
llvm/test/CodeGen/AMDGPU/nsa-vmem-hazard.mir | 2 +-
llvm/test/CodeGen/AMDGPU/occupancy-levels.ll | 13 +--
.../AMDGPU/post-ra-soft-clause-dbg-info.ll | 2 +-
llvm/test/CodeGen/AMDGPU/s_addk_i32.ll | 2 +-
llvm/test/CodeGen/AMDGPU/s_mulk_i32.ll | 2 +-
.../schedule-amdgpu-tracker-physreg-crash.ll | 4 +-
.../CodeGen/AMDGPU/soft-clause-dbg-value.mir | 4 +-
.../CodeGen/AMDGPU/spill-scavenge-offset.ll | 2 +-
llvm/test/CodeGen/AMDGPU/sram-ecc-default.ll | 12 +-
.../sramecc-subtarget-feature-disabled.ll | 5 +-
.../sramecc-subtarget-feature-enabled.ll | 5 +-
.../AMDGPU/target-id-xnack-always-on.ll | 14 +--
llvm/test/CodeGen/AMDGPU/tid-kd-xnack-off.ll | 9 +-
llvm/test/CodeGen/AMDGPU/tid-kd-xnack-on.ll | 6 +-
.../AMDGPU/tid-mul-func-xnack-all-off.ll | 3 +-
.../AMDGPU/tid-mul-func-xnack-all-on.ll | 3 +-
.../AMDGPU/tid-mul-func-xnack-any-off-1.ll | 3 +-
.../AMDGPU/tid-mul-func-xnack-any-off-2.ll | 3 +-
.../AMDGPU/tid-mul-func-xnack-any-on-1.ll | 3 +-
.../AMDGPU/tid-mul-func-xnack-any-on-2.ll | 3 +-
.../tid-mul-func-xnack-invalid-any-off-on.ll | 24 ----
.../CodeGen/AMDGPU/tid-one-func-xnack-off.ll | 3 +-
.../CodeGen/AMDGPU/tid-one-func-xnack-on.ll | 3 +-
.../xnack-subtarget-feature-disabled.ll | 7 +-
.../AMDGPU/xnack-subtarget-feature-enabled.ll | 7 +-
.../Inputs/amdgpu-sramecc-module-flag-0.ll | 6 +
.../Inputs/amdgpu-sramecc-module-flag-1.ll | 6 +
.../Inputs/amdgpu-sramecc-module-flag-any.ll | 6 +
.../Inputs/amdgpu-xnack-module-flag-0.ll | 6 +
.../Inputs/amdgpu-xnack-module-flag-1.ll | 6 +
.../Inputs/amdgpu-xnack-module-flag-any.ll | 6 +
.../Linker/amdgpu-sramecc-module-flag-0.ll | 24 ++++
.../Linker/amdgpu-sramecc-module-flag-1.ll | 23 ++++
.../Linker/amdgpu-sramecc-module-flag-any.ll | 25 +++++
.../test/Linker/amdgpu-xnack-module-flag-0.ll | 24 ++++
.../test/Linker/amdgpu-xnack-module-flag-1.ll | 23 ++++
.../Linker/amdgpu-xnack-module-flag-any.ll | 25 +++++
llvm/test/MC/AMDGPU/xnack-mask.s | 2 +-
.../Verifier/AMDGPU/module-flag-sramecc.ll | 46 ++++++++
.../test/Verifier/AMDGPU/module-flag-xnack.ll | 46 ++++++++
98 files changed, 1108 insertions(+), 357 deletions(-)
create mode 100644 clang/test/CodeGenOpenCL/amdgpu-module-flag-xnack-sramecc.cl
create mode 100644 clang/test/CodeGenOpenCL/amdgpu-xnack-any-only.cl
create mode 100644 llvm/test/CodeGen/AMDGPU/mattr-xnack-sramecc-legacy.ll
create mode 100644 llvm/test/CodeGen/AMDGPU/module-flag-sramecc.ll
create mode 100644 llvm/test/CodeGen/AMDGPU/module-flag-xnack-sramecc-combined.ll
create mode 100644 llvm/test/CodeGen/AMDGPU/module-flag-xnack.ll
delete mode 100644 llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-invalid-any-off-on.ll
create mode 100644 llvm/test/Linker/Inputs/amdgpu-sramecc-module-flag-0.ll
create mode 100644 llvm/test/Linker/Inputs/amdgpu-sramecc-module-flag-1.ll
create mode 100644 llvm/test/Linker/Inputs/amdgpu-sramecc-module-flag-any.ll
create mode 100644 llvm/test/Linker/Inputs/amdgpu-xnack-module-flag-0.ll
create mode 100644 llvm/test/Linker/Inputs/amdgpu-xnack-module-flag-1.ll
create mode 100644 llvm/test/Linker/Inputs/amdgpu-xnack-module-flag-any.ll
create mode 100644 llvm/test/Linker/amdgpu-sramecc-module-flag-0.ll
create mode 100644 llvm/test/Linker/amdgpu-sramecc-module-flag-1.ll
create mode 100644 llvm/test/Linker/amdgpu-sramecc-module-flag-any.ll
create mode 100644 llvm/test/Linker/amdgpu-xnack-module-flag-0.ll
create mode 100644 llvm/test/Linker/amdgpu-xnack-module-flag-1.ll
create mode 100644 llvm/test/Linker/amdgpu-xnack-module-flag-any.ll
create mode 100644 llvm/test/Verifier/AMDGPU/module-flag-sramecc.ll
create mode 100644 llvm/test/Verifier/AMDGPU/module-flag-xnack.ll
diff --git a/clang/include/clang/Basic/TargetOptions.h b/clang/include/clang/Basic/TargetOptions.h
index 51a2cfe14976c..466f0ef22f265 100644
--- a/clang/include/clang/Basic/TargetOptions.h
+++ b/clang/include/clang/Basic/TargetOptions.h
@@ -92,6 +92,22 @@ class TargetOptions {
/// \brief AMDGPU Printf lowering scheme
AMDGPUPrintfKind AMDGPUPrintfKindVal = AMDGPUPrintfKind::Hostcall;
+ /// \brief Enumeration values for AMDGPU xnack/sramecc settings
+ enum class AMDGPUFeatureState {
+ /// Feature state not specified and should generate most compatible code.
+ Any = 0,
+ /// Feature explicitly disabled
+ Disabled = 1,
+ /// Feature explicitly enabled
+ Enabled = 2
+ };
+
+ /// \brief AMDGPU xnack setting from -mxnack/-mno-xnack
+ AMDGPUFeatureState AMDGPUXnackState = AMDGPUFeatureState::Any;
+
+ /// \brief AMDGPU sramecc setting from -msramecc/-mno-sramecc
+ AMDGPUFeatureState AMDGPUSramEccState = AMDGPUFeatureState::Any;
+
// The code model to be used as specified by the user. Corresponds to
// CodeModel::Model enum defined in include/llvm/Support/CodeGen.h, plus
// "default" for the case when the user has not explicitly specified a
diff --git a/clang/include/clang/Options/Options.td b/clang/include/clang/Options/Options.td
index 0a24f8b35a3ea..7e2e69cc33582 100644
--- a/clang/include/clang/Options/Options.td
+++ b/clang/include/clang/Options/Options.td
@@ -5879,10 +5879,18 @@ defm cumode : SimpleMFlag<"cumode",
" execution mode (AMDGPU only)", m_amdgpu_Features_Group>;
defm tgsplit : SimpleMFlag<"tgsplit", "Enable", "Disable",
" threadgroup split execution mode (AMDGPU only)", m_amdgpu_Features_Group>;
-defm xnack : SimpleMFlag<"xnack", "Enable", "Disable",
- " XNACK (AMDGPU only)", m_amdgpu_Features_Group>;
-defm sramecc : SimpleMFlag<"sramecc", "Enable", "Disable",
- " SRAMECC (AMDGPU only)", m_amdgpu_Features_Group>;
+def mxnack : Flag<["-"], "mxnack">, Group<m_Group>,
+ Visibility<[ClangOption, CC1Option]>,
+ HelpText<"Enable XNACK (AMDGPU only)">;
+def mno_xnack : Flag<["-"], "mno-xnack">, Group<m_Group>,
+ Visibility<[ClangOption, CC1Option]>,
+ HelpText<"Disable XNACK (AMDGPU only)">;
+def msramecc : Flag<["-"], "msramecc">, Group<m_Group>,
+ Visibility<[ClangOption, CC1Option]>,
+ HelpText<"Enable SRAMECC (AMDGPU only)">;
+def mno_sramecc : Flag<["-"], "mno-sramecc">, Group<m_Group>,
+ Visibility<[ClangOption, CC1Option]>,
+ HelpText<"Disable SRAMECC (AMDGPU only)">;
defm wavefrontsize64 : SimpleMFlag<"wavefrontsize64",
"Specify wavefront size 64", "Specify wavefront size 32",
" mode (AMDGPU only)">;
diff --git a/clang/lib/Basic/Targets/AMDGPU.cpp b/clang/lib/Basic/Targets/AMDGPU.cpp
index bfa956fa9a4e3..6a87a67a688f9 100644
--- a/clang/lib/Basic/Targets/AMDGPU.cpp
+++ b/clang/lib/Basic/Targets/AMDGPU.cpp
@@ -229,6 +229,16 @@ AMDGPUTargetInfo::AMDGPUTargetInfo(const llvm::Triple &Triple,
ReadOnlyFeatures.insert(F);
}
HalfArgsAndReturns = true;
+
+ if (Opts.AMDGPUXnackState != TargetOptions::AMDGPUFeatureState::Any) {
+ OffloadArchFeatures["xnack"] =
+ Opts.AMDGPUXnackState == TargetOptions::AMDGPUFeatureState::Enabled;
+ }
+
+ if (Opts.AMDGPUSramEccState != TargetOptions::AMDGPUFeatureState::Any) {
+ OffloadArchFeatures["sramecc"] =
+ Opts.AMDGPUSramEccState == TargetOptions::AMDGPUFeatureState::Enabled;
+ }
}
void AMDGPUTargetInfo::adjust(DiagnosticsEngine &Diags, LangOptions &Opts,
diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp
index ddbebb1d70336..e095716a21f7e 100644
--- a/clang/lib/CodeGen/CodeGenModule.cpp
+++ b/clang/lib/CodeGen/CodeGenModule.cpp
@@ -1163,6 +1163,27 @@ void CodeGenModule::Release() {
getModule().addModuleFlag(llvm::Module::Error, "amdgpu_printf_kind",
MDStr);
}
+
+ const TargetOptions &TargetOpts = getTarget().getTargetOpts();
+
+ if (TargetOpts.AMDGPUXnackState != TargetOptions::AMDGPUFeatureState::Any) {
+ // TODO: Avoid emitting the xnack flag on targets which do not support
+ // xnack configuration.
+ getModule().addModuleFlag(
+ llvm::Module::Error, "amdgpu.xnack",
+ llvm::ConstantInt::get(
+ Int32Ty, TargetOpts.AMDGPUXnackState ==
+ TargetOptions::AMDGPUFeatureState::Enabled));
+ }
+
+ if (TargetOpts.AMDGPUSramEccState !=
+ TargetOptions::AMDGPUFeatureState::Any) {
+ getModule().addModuleFlag(
+ llvm::Module::Error, "amdgpu.sramecc",
+ llvm::ConstantInt::get(
+ Int32Ty, TargetOpts.AMDGPUSramEccState ==
+ TargetOptions::AMDGPUFeatureState::Enabled));
+ }
}
// Emit a global array containing all external kernels or device variables
diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp
index e3288c81d4c95..d58b077c85acd 100644
--- a/clang/lib/Driver/ToolChains/Clang.cpp
+++ b/clang/lib/Driver/ToolChains/Clang.cpp
@@ -1535,6 +1535,15 @@ void Clang::AddARMTargetArgs(const llvm::Triple &Triple, const ArgList &Args,
AddUnalignedAccessWarning(CmdArgs);
}
+void Clang::AddAMDGPUTargetArgs(const ArgList &Args,
+ ArgStringList &CmdArgs) const {
+ // Pass through -mxnack/-mno-xnack and -msramecc/-mno-sramecc flags to cc1.
+ if (Arg *A = Args.getLastArg(options::OPT_mxnack, options::OPT_mno_xnack))
+ A->render(Args, CmdArgs);
+ if (Arg *A = Args.getLastArg(options::OPT_msramecc, options::OPT_mno_sramecc))
+ A->render(Args, CmdArgs);
+}
+
void Clang::RenderTargetOptions(const llvm::Triple &EffectiveTriple,
const ArgList &Args, bool KernelOrKext,
ArgStringList &CmdArgs) const {
@@ -1562,6 +1571,10 @@ void Clang::RenderTargetOptions(const llvm::Triple &EffectiveTriple,
AddAArch64TargetArgs(Args, CmdArgs);
break;
+ case llvm::Triple::amdgcn:
+ AddAMDGPUTargetArgs(Args, CmdArgs);
+ break;
+
case llvm::Triple::loongarch32:
case llvm::Triple::loongarch64:
AddLoongArchTargetArgs(Args, CmdArgs);
diff --git a/clang/lib/Driver/ToolChains/Clang.h b/clang/lib/Driver/ToolChains/Clang.h
index 9adad5c5430f2..23f270ffb76a3 100644
--- a/clang/lib/Driver/ToolChains/Clang.h
+++ b/clang/lib/Driver/ToolChains/Clang.h
@@ -51,6 +51,8 @@ class LLVM_LIBRARY_VISIBILITY Clang : public Tool {
void AddAArch64TargetArgs(const llvm::opt::ArgList &Args,
llvm::opt::ArgStringList &CmdArgs) const;
+ void AddAMDGPUTargetArgs(const llvm::opt::ArgList &Args,
+ llvm::opt::ArgStringList &CmdArgs) const;
void AddARMTargetArgs(const llvm::Triple &Triple,
const llvm::opt::ArgList &Args,
llvm::opt::ArgStringList &CmdArgs,
diff --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp
index d2847739e3143..e8f9bc5a60737 100644
--- a/clang/lib/Frontend/CompilerInvocation.cpp
+++ b/clang/lib/Frontend/CompilerInvocation.cpp
@@ -5024,6 +5024,18 @@ static void GenerateTargetArgs(const TargetOptions &Opts,
if (!Opts.DarwinTargetVariantSDKVersion.empty())
GenerateArg(Consumer, OPT_darwin_target_variant_sdk_version_EQ,
Opts.DarwinTargetVariantSDKVersion.getAsString());
+
+ // Generate AMDGPU xnack and sramecc flags.
+ if (Opts.AMDGPUXnackState == TargetOptions::AMDGPUFeatureState::Enabled)
+ GenerateArg(Consumer, OPT_mxnack);
+ else if (Opts.AMDGPUXnackState == TargetOptions::AMDGPUFeatureState::Disabled)
+ GenerateArg(Consumer, OPT_mno_xnack);
+
+ if (Opts.AMDGPUSramEccState == TargetOptions::AMDGPUFeatureState::Enabled)
+ GenerateArg(Consumer, OPT_msramecc);
+ else if (Opts.AMDGPUSramEccState ==
+ TargetOptions::AMDGPUFeatureState::Disabled)
+ GenerateArg(Consumer, OPT_mno_sramecc);
}
static bool ParseTargetArgs(TargetOptions &Opts, ArgList &Args,
@@ -5055,6 +5067,21 @@ static bool ParseTargetArgs(TargetOptions &Opts, ArgList &Args,
Opts.DarwinTargetVariantSDKVersion = Version;
}
+ if (Arg *A = Args.getLastArg(options::OPT_mxnack, options::OPT_mno_xnack)) {
+ bool IsEnabled = A->getOption().matches(options::OPT_mxnack);
+ Opts.AMDGPUXnackState = IsEnabled
+ ? TargetOptions::AMDGPUFeatureState::Enabled
+ : TargetOptions::AMDGPUFeatureState::Disabled;
+ }
+
+ if (Arg *A =
+ Args.getLastArg(options::OPT_msramecc, options::OPT_mno_sramecc)) {
+ bool IsEnabled = A->getOption().matches(options::OPT_msramecc);
+ Opts.AMDGPUSramEccState = IsEnabled
+ ? TargetOptions::AMDGPUFeatureState::Enabled
+ : TargetOptions::AMDGPUFeatureState::Disabled;
+ }
+
return Diags.getNumErrors() == NumErrorsBefore;
}
diff --git a/clang/test/CIR/CodeGenHIP/target-features.hip b/clang/test/CIR/CodeGenHIP/target-features.hip
index 8d414edcd8e2c..afce90caca435 100644
--- a/clang/test/CIR/CodeGenHIP/target-features.hip
+++ b/clang/test/CIR/CodeGenHIP/target-features.hip
@@ -18,17 +18,17 @@
// only the delta (differing features) is emitted on cir.target-features.
// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fclangir \
-// RUN: -fcuda-is-device -target-cpu gfx900 -target-feature +xnack \
+// RUN: -fcuda-is-device -target-cpu gfx900 -mxnack \
// RUN: -emit-cir %s -o %t-delta.cir
// RUN: FileCheck --check-prefix=CIR-DELTA %s --input-file=%t-delta.cir
// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fclangir \
-// RUN: -fcuda-is-device -target-cpu gfx900 -target-feature +xnack \
+// RUN: -fcuda-is-device -target-cpu gfx900 -mxnack \
// RUN: -emit-llvm %s -o %t-delta.cir.ll
// RUN: FileCheck --check-prefix=LLVM-DELTA %s --input-file=%t-delta.cir.ll
// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip \
-// RUN: -fcuda-is-device -target-cpu gfx900 -target-feature +xnack \
+// RUN: -fcuda-is-device -target-cpu gfx900 -mxnack \
// RUN: -emit-llvm %s -o %t-delta.ll
// RUN: FileCheck --check-prefix=LLVM-DELTA %s --input-file=%t-delta.ll
@@ -49,17 +49,18 @@ __device__ void device_fn() {}
// LLVM-DAG: attributes #[[K_ATTR]] = {{.*}}"target-cpu"="gfx900"
// LLVM-DAG: attributes #[[D_ATTR]] = {{.*}}"target-cpu"="gfx900"
-// AMDGPU with gfx900 + an extra +xnack feature: only the delta is emitted.
+// AMDGPU with gfx900 + xnack enabled via -mxnack: emitted as module flag.
// CIR-DELTA: cir.func{{.*}} @_Z6kernelv()
// CIR-DELTA-SAME: "cir.target-cpu" = "gfx900"
-// CIR-DELTA-SAME: "cir.target-features" = "+xnack"
+// CIR-DELTA-NOT: cir.target-features
// CIR-DELTA: cir.func{{.*}} @_Z9device_fnv()
// CIR-DELTA-SAME: "cir.target-cpu" = "gfx900"
-// CIR-DELTA-SAME: "cir.target-features" = "+xnack"
+// CIR-DELTA-NOT: cir.target-features
// LLVM-DELTA: define{{.*}} void @_Z6kernelv(){{.*}} #[[K_ATTR_D:[0-9]+]]
// LLVM-DELTA: define{{.*}} void @_Z9device_fnv(){{.*}} #[[D_ATTR_D:[0-9]+]]
-// LLVM-DELTA-DAG: attributes #[[K_ATTR_D]] = {{.*}}"target-cpu"="gfx900"{{.*}}"target-features"="+xnack"
-// LLVM-DELTA-DAG: attributes #[[D_ATTR_D]] = {{.*}}"target-cpu"="gfx900"{{.*}}"target-features"="+xnack"
+// LLVM-DELTA-DAG: attributes #[[K_ATTR_D]] = {{.*}}"target-cpu"="gfx900"
+// LLVM-DELTA-DAG: attributes #[[D_ATTR_D]] = {{.*}}"target-cpu"="gfx900"
+// LLVM-DELTA-DAG: !{i32 1, !"amdgpu.xnack", i32 1}
diff --git a/clang/test/CodeGenCXX/dynamic-cast-address-space.cpp b/clang/test/CodeGenCXX/dynamic-cast-address-space.cpp
index e5873c1a8de90..2d48708601ccf 100644
--- a/clang/test/CodeGenCXX/dynamic-cast-address-space.cpp
+++ b/clang/test/CodeGenCXX/dynamic-cast-address-space.cpp
@@ -107,9 +107,9 @@ const B& f(A *a) {
// CHECK: attributes #[[ATTR3]] = { nounwind }
// CHECK: attributes #[[ATTR4]] = { noreturn }
//.
-// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR0]] = { mustprogress noinline optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+16-bit-insts,+add-min-max-insts,+ashr-pk-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-fmin-fmax-global-f32,+atomic-fmin-fmax-global-f64,+atomic-global-pk-add-bf16-inst,+bf16-cvt-insts,+bf16-pk-insts,+bf16-trans-insts,+bf8-cvt-scale-insts,+bitop3-insts,+ci-insts,+clusters,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot13-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+f16bf16-to-fp6bf6-cvt-scale-insts,+f32-to-f16bf16-cvt-sr-insts,+flat-global-insts,+fp4-cvt-scale-insts,+fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-cvt-scale-insts,+fp8-insts,+fp8e5m3-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx1251-gemm-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gfx950-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mcast-load-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+permlane16-swap,+permlane32-swap,+pk-add-min-max-insts,+prng-inst,+qsad-insts,+s-memrealtime,+s-memtime-inst,+s-wakeup-barrier-inst,+sad-insts,+setprio-inc-wg-inst,+swmmac-gfx1200-insts,+swmmac-gfx1250-insts,+tanh-insts,+tensor-cvt-lut-insts,+transpose-load-f4f6-insts,+vmem-pref-insts,+vmem-to-lds-load-insts,+wavefrontsize32,+wavefrontsize64,+wmma-128b-insts,+wmma-256b-insts,+xf32-insts" }
+// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR0]] = { mustprogress noinline optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+16-bit-insts,+add-min-max-insts,+ashr-pk-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-fmin-fmax-global-f32,+atomic-fmin-fmax-global-f64,+atomic-global-pk-add-bf16-inst,+bf16-cvt-insts,+bf16-pk-insts,+bf16-trans-insts,+bf8-cvt-scale-insts,+bitop3-insts,+ci-insts,+clusters,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot13-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+f16bf16-to-fp6bf6-cvt-scale-insts,+f32-to-f16bf16-cvt-sr-insts,+flat-global-insts,+fp4-cvt-scale-insts,+fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-cvt-scale-insts,+fp8-insts,+fp8e5m3-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx1251-gemm-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gfx950-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mcast-load-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+permlane16-swap,+permlane32-swap,+pk-add-min-max-insts,+prng-inst,+qsad-insts,+s-memrealtime,+s-memtime-inst,+s-wakeup-barrier-inst,+sad-insts,+setprio-inc-wg-inst,+swmmac-gfx1200-insts,+swmmac-gfx1250-insts,+tanh-insts,+tensor-cvt-lut-insts,+transpose-load-f4f6-insts,+vmem-pref-insts,+vmem-to-lds-load-insts,+wavefrontsize32,+wavefrontsize64,+wmma-128b-insts,+wmma-256b-insts,+xf32-insts,+xnack-any-only" }
// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR1:[0-9]+]] = { nounwind willreturn memory(read) }
-// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR2:[0-9]+]] = { "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+16-bit-insts,+add-min-max-insts,+ashr-pk-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-fmin-fmax-global-f32,+atomic-fmin-fmax-global-f64,+atomic-global-pk-add-bf16-inst,+bf16-cvt-insts,+bf16-pk-insts,+bf16-trans-insts,+bf8-cvt-scale-insts,+bitop3-insts,+ci-insts,+clusters,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot13-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+f16bf16-to-fp6bf6-cvt-scale-insts,+f32-to-f16bf16-cvt-sr-insts,+flat-global-insts,+fp4-cvt-scale-insts,+fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-cvt-scale-insts,+fp8-insts,+fp8e5m3-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx1251-gemm-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gfx950-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mcast-load-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+permlane16-swap,+permlane32-swap,+pk-add-min-max-insts,+prng-inst,+qsad-insts,+s-memrealtime,+s-memtime-inst,+s-wakeup-barrier-inst,+sad-insts,+setprio-inc-wg-inst,+swmmac-gfx1200-insts,+swmmac-gfx1250-insts,+tanh-insts,+tensor-cvt-lut-insts,+transpose-load-f4f6-insts,+vmem-pref-insts,+vmem-to-lds-load-insts,+wavefrontsize32,+wavefrontsize64,+wmma-128b-insts,+wmma-256b-insts,+xf32-insts" }
+// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR2:[0-9]+]] = { "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+16-bit-insts,+add-min-max-insts,+ashr-pk-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-fmin-fmax-global-f32,+atomic-fmin-fmax-global-f64,+atomic-global-pk-add-bf16-inst,+bf16-cvt-insts,+bf16-pk-insts,+bf16-trans-insts,+bf8-cvt-scale-insts,+bitop3-insts,+ci-insts,+clusters,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot13-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+f16bf16-to-fp6bf6-cvt-scale-insts,+f32-to-f16bf16-cvt-sr-insts,+flat-global-insts,+fp4-cvt-scale-insts,+fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-cvt-scale-insts,+fp8-insts,+fp8e5m3-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx1251-gemm-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gfx950-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mcast-load-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+permlane16-swap,+permlane32-swap,+pk-add-min-max-insts,+prng-inst,+qsad-insts,+s-memrealtime,+s-memtime-inst,+s-wakeup-barrier-inst,+sad-insts,+setprio-inc-wg-inst,+swmmac-gfx1200-insts,+swmmac-gfx1250-insts,+tanh-insts,+tensor-cvt-lut-insts,+transpose-load-f4f6-insts,+vmem-pref-insts,+vmem-to-lds-load-insts,+wavefrontsize32,+wavefrontsize64,+wmma-128b-insts,+wmma-256b-insts,+xf32-insts,+xnack-any-only" }
// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR3]] = { nounwind }
// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR4]] = { noreturn }
//.
diff --git a/clang/test/CodeGenOpenCL/amdgpu-module-flag-xnack-sramecc.cl b/clang/test/CodeGenOpenCL/amdgpu-module-flag-xnack-sramecc.cl
new file mode 100644
index 0000000000000..439a815f5c42d
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/amdgpu-module-flag-xnack-sramecc.cl
@@ -0,0 +1,22 @@
+// Test that xnack and sramecc module flags are emitted based on -m flags
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx90a \
+// RUN: -mxnack -mno-sramecc \
+// RUN: -emit-llvm -o - %s | FileCheck %s --check-prefixes=XNACK-ON,SRAMECC-OFF
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx90a \
+// RUN: -mno-xnack -msramecc \
+// RUN: -emit-llvm -o - %s | FileCheck %s --check-prefixes=XNACK-OFF,SRAMECC-ON
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx90a \
+// RUN: -emit-llvm -o - %s | FileCheck %s --check-prefix=NO-FLAGS
+
+// XNACK-ON-DAG: !{i32 1, !"amdgpu.xnack", i32 1}
+// XNACK-OFF-DAG: !{i32 1, !"amdgpu.xnack", i32 0}
+// SRAMECC-ON-DAG: !{i32 1, !"amdgpu.sramecc", i32 1}
+// SRAMECC-OFF-DAG: !{i32 1, !"amdgpu.sramecc", i32 0}
+
+// When no explicit xnack/sramecc feature is set, no module flags are emitted
+// NO-FLAGS-NOT: !"amdgpu.xnack"
+// NO-FLAGS-NOT: !"amdgpu.sramecc"
+
+__attribute__((device)) void test() {}
diff --git a/clang/test/CodeGenOpenCL/amdgpu-xnack-any-only.cl b/clang/test/CodeGenOpenCL/amdgpu-xnack-any-only.cl
new file mode 100644
index 0000000000000..709ca1acf4cb3
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/amdgpu-xnack-any-only.cl
@@ -0,0 +1,27 @@
+// Test that xnack module flags are emitted for all targets, regardless of support.
+// Targets without FEATURE_XNACK_ON_OFF_MODES (like gfx12-5-generic, gfx1250, gfx1251)
+// will ignore the module flag during codegen, but it is still emitted by clang.
+// TODO: In the future, clang should not emit the flag for targets that don't support
+// xnack control.
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx12-5-generic \
+// RUN: -mxnack -emit-llvm -o - %s | FileCheck %s --check-prefix=XNACK-ON
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx12-5-generic \
+// RUN: -mno-xnack -emit-llvm -o - %s | FileCheck %s --check-prefix=XNACK-OFF
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx12-5-generic \
+// RUN: -emit-llvm -o - %s | FileCheck %s --check-prefix=NO-FLAGS
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx1250 \
+// RUN: -mxnack -emit-llvm -o - %s | FileCheck %s --check-prefix=XNACK-ON
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx1250 \
+// RUN: -mno-xnack -emit-llvm -o - %s | FileCheck %s --check-prefix=XNACK-OFF
+
+// Module flags are emitted regardless of target support
+// XNACK-ON-DAG: !{i32 1, !"amdgpu.xnack", i32 1}
+// XNACK-OFF-DAG: !{i32 1, !"amdgpu.xnack", i32 0}
+// NO-FLAGS-NOT: !"amdgpu.xnack"
+
+kernel void test() {}
diff --git a/clang/test/Driver/amdgpu-features.c b/clang/test/Driver/amdgpu-features.c
index c756b91379180..9513f05fbb58b 100644
--- a/clang/test/Driver/amdgpu-features.c
+++ b/clang/test/Driver/amdgpu-features.c
@@ -1,14 +1,14 @@
// RUN: %clang -### --target=amdgcn-amdhsa -mcpu=gfx900:xnack+ -nogpulib %s 2>&1 | FileCheck --check-prefix=XNACK %s
-// XNACK: "-target-feature" "+xnack"
+// XNACK: "-mxnack"
// RUN: %clang -### -target amdgcn-amdpal -mcpu=gfx900:xnack- %s 2>&1 | FileCheck --check-prefix=NO-XNACK %s
-// NO-XNACK: "-target-feature" "-xnack"
+// NO-XNACK: "-mno-xnack"
// RUN: %clang -### -target amdgcn-mesa3d -mcpu=gfx908:sramecc+ %s 2>&1 | FileCheck --check-prefix=SRAM-ECC %s
-// SRAM-ECC: "-target-feature" "+sramecc"
+// SRAM-ECC: "-msramecc"
// RUN: %clang -### --target=amdgcn-amdhsa -mcpu=gfx908:sramecc- -nogpulib %s 2>&1 | FileCheck --check-prefix=NO-SRAM-ECC %s
-// NO-SRAM-ECC: "-target-feature" "-sramecc"
+// NO-SRAM-ECC: "-mno-sramecc"
// RUN: %clang -### -target amdgcn -mcpu=gfx90a -mtgsplit %s 2>&1 | FileCheck --check-prefix=TGSPLIT %s
// RUN: %clang -### -target amdgcn -mcpu=gfx90a -mno-tgsplit %s 2>&1 | FileCheck --check-prefix=NO-TGSPLIT %s
diff --git a/clang/test/Driver/amdgpu-openmp-toolchain.c b/clang/test/Driver/amdgpu-openmp-toolchain.c
index 4de585e7c6238..49671710133ea 100644
--- a/clang/test/Driver/amdgpu-openmp-toolchain.c
+++ b/clang/test/Driver/amdgpu-openmp-toolchain.c
@@ -63,7 +63,7 @@
// RUN: %clang -### -target x86_64-pc-linux-gnu -fopenmp --offload-arch=gfx90a:sramecc-:xnack+ \
// RUN: -nogpulib %s 2>&1 | FileCheck %s --check-prefix=CHECK-TARGET-ID
-// CHECK-TARGET-ID: "-cc1" "-triple" "amdgcn-amd-amdhsa" {{.*}} "-target-cpu" "gfx90a" "-target-feature" "+xnack" "-target-feature" "-sramecc"
+// CHECK-TARGET-ID: "-cc1" "-triple" "amdgcn-amd-amdhsa" {{.*}} "-target-cpu" "gfx90a" "-mxnack" "-mno-sramecc"
// CHECK-TARGET-ID: llvm-offload-binary{{.*}}arch=gfx90a:sramecc-:xnack+,kind=openmp
// RUN: not %clang -### -target x86_64-pc-linux-gnu -fopenmp --offload-arch=gfx90a,gfx90a:xnack+ \
diff --git a/clang/test/Driver/amdgpu-toolchain.c b/clang/test/Driver/amdgpu-toolchain.c
index 135129b739603..25f9a9094aa77 100644
--- a/clang/test/Driver/amdgpu-toolchain.c
+++ b/clang/test/Driver/amdgpu-toolchain.c
@@ -25,11 +25,11 @@
// RUN: %clang -### --target=amdgcn-amd-amdhsa -mcpu=gfx90a:xnack+:sramecc- -nogpulib \
// RUN: -L. -flto -fconvergent-functions %s 2>&1 | FileCheck -check-prefix=LTO %s
// LTO: clang{{.*}}"-flto=full"{{.*}}"-fconvergent-functions"
-// LTO: ld.lld{{.*}}"-plugin-opt=mcpu=gfx90a"{{.*}}"-plugin-opt=-mattr=+xnack,-sramecc"{{.*}}
+// LTO: ld.lld{{.*}}"-plugin-opt=mcpu=gfx90a"
// RUN: %clang -### --target=amdgcn-amd-amdhsa -mcpu=gfx90a:xnack+:sramecc- -nogpulib \
// RUN: -L. -fconvergent-functions %s 2>&1 | FileCheck -check-prefix=MCPU %s
-// MCPU: ld.lld{{.*}}"-plugin-opt=mcpu=gfx90a"{{.*}}"-plugin-opt=-mattr=+xnack,-sramecc"{{.*}}
+// MCPU: ld.lld{{.*}}"-plugin-opt=mcpu=gfx90a"
// RUN: %clang -### --target=amdgcn-amd-amdhsa -mcpu=gfx906 -nogpulib \
// RUN: -fuse-ld=ld %s 2>&1 | FileCheck -check-prefixes=LD %s
diff --git a/clang/test/Driver/amdgpu-xnack-sramecc-flags.c b/clang/test/Driver/amdgpu-xnack-sramecc-flags.c
index c40dabcd4f645..58b3c2c6ba612 100644
--- a/clang/test/Driver/amdgpu-xnack-sramecc-flags.c
+++ b/clang/test/Driver/amdgpu-xnack-sramecc-flags.c
@@ -1,68 +1,71 @@
-// Test for -mxnack/-mno-xnack and -msramecc/-mno-sramecc flags
+// Test for -mxnack/-mno-xnack and -msramecc/-mno-sramecc flags, which should be
+// forwarded to cc1.
// RUN: %clang -### --target=amdgcn-amd-amdhsa -mcpu=gfx90a -mxnack %s 2>&1 | \
// RUN: FileCheck -check-prefix=XNACK-ON %s
-// XNACK-ON: "-target-feature" "+xnack"
+// XNACK-ON: "-mxnack"
// RUN: %clang -### --target=amdgcn-amd-amdhsa -mcpu=gfx90a -mno-xnack %s 2>&1 | \
// RUN: FileCheck -check-prefix=XNACK-OFF %s
-// XNACK-OFF: "-target-feature" "-xnack"
+// XNACK-OFF: "-mno-xnack"
// RUN: %clang -### --target=amdgcn-amd-amdhsa -mcpu=gfx90a -msramecc %s 2>&1 | \
// RUN: FileCheck -check-prefix=SRAMECC-ON %s
-// SRAMECC-ON: "-target-feature" "+sramecc"
+// SRAMECC-ON: "-msramecc"
// RUN: %clang -### --target=amdgcn-amd-amdhsa -mcpu=gfx90a -mno-sramecc %s 2>&1 | \
// RUN: FileCheck -check-prefix=SRAMECC-OFF %s
-// SRAMECC-OFF: "-target-feature" "-sramecc"
+// SRAMECC-OFF: "-mno-sramecc"
// Test that target ID takes precedence over explicit flags
// RUN: %clang -### --target=amdgcn-amd-amdhsa -mcpu=gfx90a:xnack+ -mno-xnack %s 2>&1 | \
// RUN: FileCheck -check-prefix=TARGETID-OVERRIDES-XNACK %s
-// TARGETID-OVERRIDES-XNACK: "-target-feature" "+xnack"
+// TARGETID-OVERRIDES-XNACK: "-mxnack"
// RUN: %clang -### --target=amdgcn-amd-amdhsa -mcpu=gfx90a:xnack- -mxnack %s 2>&1 | \
// RUN: FileCheck -check-prefix=TARGETID-OVERRIDES-XNACK-OFF %s
-// TARGETID-OVERRIDES-XNACK-OFF: "-target-feature" "-xnack"
+// TARGETID-OVERRIDES-XNACK-OFF: "-mno-xnack"
// RUN: %clang -### --target=amdgcn-amd-amdhsa -mcpu=gfx90a:sramecc+ -mno-sramecc %s 2>&1 | \
// RUN: FileCheck -check-prefix=TARGETID-OVERRIDES-SRAMECC %s
-// TARGETID-OVERRIDES-SRAMECC: "-target-feature" "+sramecc"
+// TARGETID-OVERRIDES-SRAMECC: "-msramecc"
// RUN: %clang -### --target=amdgcn-amd-amdhsa -mcpu=gfx90a:sramecc- -msramecc %s 2>&1 | \
// RUN: FileCheck -check-prefix=TARGETID-OVERRIDES-SRAMECC-OFF %s
-// TARGETID-OVERRIDES-SRAMECC-OFF: "-target-feature" "-sramecc"
+// TARGETID-OVERRIDES-SRAMECC-OFF: "-mno-sramecc"
// Test combining both flags
// RUN: %clang -### --target=amdgcn-amd-amdhsa -mcpu=gfx90a -mxnack -msramecc %s 2>&1 | \
// RUN: FileCheck -check-prefixes=BOTH-ON %s
-// BOTH-ON: "-target-feature" "+xnack"
-// BOTH-ON-SAME: "-target-feature" "+sramecc"
+// BOTH-ON: "-mxnack"
+// BOTH-ON-SAME: "-msramecc"
// RUN: %clang -### --target=amdgcn-amd-amdhsa -mcpu=gfx90a -mno-xnack -mno-sramecc %s 2>&1 | \
// RUN: FileCheck -check-prefixes=BOTH-OFF %s
-// BOTH-OFF: "-target-feature" "-xnack"
-// BOTH-OFF-SAME: "-target-feature" "-sramecc"
+// BOTH-OFF: "-mno-xnack"
+// BOTH-OFF-SAME: "-mno-sramecc"
// Test that target ID without explicit features doesn't synthesize flags
// RUN: %clang -### --target=amdgcn-amd-amdhsa -mcpu=gfx90a %s 2>&1 | \
// RUN: FileCheck -check-prefix=NO-FEATURES %s
-// NO-FEATURES-NOT: "-target-feature" "{{[+-]}}xnack"
-// NO-FEATURES-NOT: "-target-feature" "{{[+-]}}sramecc"
+// NO-FEATURES-NOT: "-mxnack"
+// NO-FEATURES-NOT: "-mno-xnack"
+// NO-FEATURES-NOT: "-msramecc"
+// NO-FEATURES-NOT: "-mno-sramecc"
-// Test target ID features are synthesized
+// Test target ID features are synthesized as flags
// RUN: %clang -### --target=amdgcn-amd-amdhsa -mcpu=gfx90a:xnack+ %s 2>&1 | \
// RUN: FileCheck -check-prefix=TARGETID-XNACK %s
-// TARGETID-XNACK: "-target-feature" "+xnack"
+// TARGETID-XNACK: "-mxnack"
// RUN: %clang -### --target=amdgcn-amd-amdhsa -mcpu=gfx90a:sramecc+ %s 2>&1 | \
// RUN: FileCheck -check-prefix=TARGETID-SRAMECC %s
-// TARGETID-SRAMECC: "-target-feature" "+sramecc"
+// TARGETID-SRAMECC: "-msramecc"
// RUN: %clang -### --target=amdgcn-amd-amdhsa -mcpu=gfx90a:xnack+:sramecc+ %s 2>&1 | \
// RUN: FileCheck -check-prefixes=TARGETID-BOTH %s
-// TARGETID-BOTH: "-target-feature" "+xnack"
-// TARGETID-BOTH-SAME: "-target-feature" "+sramecc"
+// TARGETID-BOTH: "-mxnack"
+// TARGETID-BOTH-SAME: "-msramecc"
//
// Offload tests
@@ -72,16 +75,16 @@
// RUN: %clang -### --target=x86_64-unknown-linux-gnu -fopenmp --offload-arch=gfx90a:xnack+:sramecc- \
// RUN: -nogpulib %s 2>&1 | FileCheck -check-prefix=OMP-TARGETID %s
// OMP-TARGETID: "-cc1" "-triple" "amdgcn-amd-amdhsa" {{.*}} "-target-cpu" "gfx90a"
-// OMP-TARGETID-SAME: "-target-feature" "+xnack"
-// OMP-TARGETID-SAME: "-target-feature" "-sramecc"
+// OMP-TARGETID-SAME: "-mxnack"
+// OMP-TARGETID-SAME: "-mno-sramecc"
// Test offload using -fopenmp-targets with target ID
// RUN: %clang -### --target=x86_64-unknown-linux-gnu -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa \
// RUN: -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx908:xnack-:sramecc+ \
// RUN: -nogpulib %s 2>&1 | FileCheck -check-prefix=OMP-MARCH %s
// OMP-MARCH: "-cc1" "-triple" "amdgcn-amd-amdhsa" {{.*}} "-target-cpu" "gfx908"
-// OMP-MARCH-SAME: "-target-feature" "-xnack"
-// OMP-MARCH-SAME: "-target-feature" "+sramecc"
+// OMP-MARCH-SAME: "-mno-xnack"
+// OMP-MARCH-SAME: "-msramecc"
// Test offload with explicit device flags using -Xopenmp-target
// RUN: %clang -### --target=x86_64-unknown-linux-gnu -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa \
@@ -90,8 +93,8 @@
// RUN: -Xopenmp-target=amdgcn-amd-amdhsa -mno-sramecc \
// RUN: -nogpulib %s 2>&1 | FileCheck -check-prefix=OMP-FLAGS %s
// OMP-FLAGS: "-cc1" "-triple" "amdgcn-amd-amdhsa" {{.*}} "-target-cpu" "gfx90a"
-// OMP-FLAGS-SAME: "-target-feature" "+xnack"
-// OMP-FLAGS-SAME: "-target-feature" "-sramecc"
+// OMP-FLAGS-SAME: "-mxnack"
+// OMP-FLAGS-SAME: "-mno-sramecc"
// Test offload with target ID taking precedence over explicit flags
// RUN: %clang -### --target=x86_64-unknown-linux-gnu -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa \
@@ -99,7 +102,7 @@
// RUN: -Xopenmp-target=amdgcn-amd-amdhsa -mxnack \
// RUN: -nogpulib %s 2>&1 | FileCheck -check-prefix=OMP-TARGETID-WINS %s
// OMP-TARGETID-WINS: "-cc1" "-triple" "amdgcn-amd-amdhsa" {{.*}} "-target-cpu" "gfx90a"
-// OMP-TARGETID-WINS-SAME: "-target-feature" "-xnack"
+// OMP-TARGETID-WINS-SAME: "-mno-xnack"
// Test offload using base architecture gfx90a with -mxnack flag for xnack+
// RUN: %clang -### --target=x86_64-unknown-linux-gnu -fopenmp \
@@ -107,7 +110,7 @@
// RUN: -Xopenmp-target=amdgcn-amd-amdhsa -mxnack \
// RUN: -nogpulib %s 2>&1 | FileCheck -check-prefix=OMP-GFX90A-XNACK-ON %s
// OMP-GFX90A-XNACK-ON: "-cc1" "-triple" "amdgcn-amd-amdhsa" {{.*}} "-target-cpu" "gfx90a"
-// OMP-GFX90A-XNACK-ON-SAME: "-target-feature" "+xnack"
+// OMP-GFX90A-XNACK-ON-SAME: "-mxnack"
// Test offload using base architecture gfx90a with -mno-xnack flag for xnack-
// RUN: %clang -### --target=x86_64-unknown-linux-gnu -fopenmp \
@@ -115,7 +118,7 @@
// RUN: -Xopenmp-target=amdgcn-amd-amdhsa -mno-xnack \
// RUN: -nogpulib %s 2>&1 | FileCheck -check-prefix=OMP-GFX90A-XNACK-OFF %s
// OMP-GFX90A-XNACK-OFF: "-cc1" "-triple" "amdgcn-amd-amdhsa" {{.*}} "-target-cpu" "gfx90a"
-// OMP-GFX90A-XNACK-OFF-SAME: "-target-feature" "-xnack"
+// OMP-GFX90A-XNACK-OFF-SAME: "-mno-xnack"
// Test offload with multiple device compilations for same base architecture.
// To get both xnack+ and xnack- for gfx90a in the same invocation, you must use
@@ -124,9 +127,9 @@
// RUN: --offload-arch=gfx90a:xnack+ --offload-arch=gfx90a:xnack- -mxnack \
// RUN: -nogpulib %s 2>&1 | FileCheck -check-prefix=OMP-MULTI-XNACK %s
// OMP-MULTI-XNACK: "-cc1" "-triple" "amdgcn-amd-amdhsa" {{.*}} "-target-cpu" "gfx90a"
-// OMP-MULTI-XNACK-SAME: "-target-feature" "+xnack"
+// OMP-MULTI-XNACK-SAME: "-mxnack"
// OMP-MULTI-XNACK: "-cc1" "-triple" "amdgcn-amd-amdhsa" {{.*}} "-target-cpu" "gfx90a"
-// OMP-MULTI-XNACK-SAME: "-target-feature" "-xnack"
+// OMP-MULTI-XNACK-SAME: "-mno-xnack"
// Test that -Xopenmp-target flags apply to all targets with matching triple.
// When compiling for multiple different base architectures (gfx906, gfx90a),
@@ -136,9 +139,9 @@
// RUN: -Xopenmp-target=amdgcn-amd-amdhsa -mxnack \
// RUN: -nogpulib %s 2>&1 | FileCheck -check-prefix=OMP-MULTI-ARCH %s
// OMP-MULTI-ARCH: "-cc1" "-triple" "amdgcn-amd-amdhsa" {{.*}} "-target-cpu" "gfx906"
-// OMP-MULTI-ARCH-SAME: "-target-feature" "+xnack"
+// OMP-MULTI-ARCH-SAME: "-mxnack"
// OMP-MULTI-ARCH: "-cc1" "-triple" "amdgcn-amd-amdhsa" {{.*}} "-target-cpu" "gfx90a"
-// OMP-MULTI-ARCH-SAME: "-target-feature" "+xnack"
+// OMP-MULTI-ARCH-SAME: "-mxnack"
// Test that top-level -mxnack flags (not specified to the device are ignored).
// TODO: Should this be forwarded?
diff --git a/clang/test/Driver/hip-sanitize-options.hip b/clang/test/Driver/hip-sanitize-options.hip
index 16eccf4a76013..964d89a609bb9 100644
--- a/clang/test/Driver/hip-sanitize-options.hip
+++ b/clang/test/Driver/hip-sanitize-options.hip
@@ -118,41 +118,41 @@
// XNACK: warning: ignoring 'leak' in '-fsanitize=leak' option as it is not currently supported for target 'amdgcn-amd-amdhsa'
// XNACK: warning: ignoring '-fsanitize=address' option for offload arch 'gfx900:xnack-' as it is not currently supported there. Use it with an offload arch containing 'xnack+' instead
// XNACK: warning: ignoring '-fsanitize=address' option for offload arch 'gfx906' as it is not currently supported there. Use it with an offload arch containing 'xnack+' instead
-// XNACK-DAG: {{"[^"]*clang[^"]*".* "-mlink-bitcode-file" ".*asanrtl.bc".* "-target-cpu" "gfx900".* "-target-feature" "\+xnack".* "-fsanitize=address"}}
-// XNACK-DAG: {{"[^"]*clang[^"]*".* "-target-cpu" "gfx900".* "-target-feature" "-xnack"}}
+// XNACK-DAG: {{"[^"]*clang[^"]*".* "-mlink-bitcode-file" ".*asanrtl.bc".* "-target-cpu" "gfx900".* "-mxnack".* "-fsanitize=address"}}
+// XNACK-DAG: {{"[^"]*clang[^"]*".* "-target-cpu" "gfx900".* "-mno-xnack"}}
// XNACK-DAG: {{"[^"]*clang[^"]*".* "-target-cpu" "gfx906"}}
// XNACK-DAG: {{"[^"]*clang[^"]*".* "-triple" "x86_64-unknown-linux-gnu".* "-fsanitize=address,leak"}}
-// XNACKNEG-NOT: {{"[^"]*clang[^"]*".* "-target-cpu" "gfx900".* "-target-feature" "+xnack".* "-fsanitize=address,leak"}}
-// XNACKNEG-NOT: {{"[^"]*clang[^"]*".* "-target-cpu" "gfx900".* "-target-feature" "-xnack".* "-fsanitize=address,leak"}}
+// XNACKNEG-NOT: {{"[^"]*clang[^"]*".* "-target-cpu" "gfx900".* "-mxnack".* "-fsanitize=address,leak"}}
+// XNACKNEG-NOT: {{"[^"]*clang[^"]*".* "-target-cpu" "gfx900".* "-mno-xnack".* "-fsanitize=address,leak"}}
// XNACKNEG-NOT: {{"[^"]*clang[^"]*".* "-target-cpu" "gfx906".* "-fsanitize=address,leak"}}
-// XNACKNEG-NOT: {{"[^"]*clang[^"]*".* "-target-cpu" "gfx900".* "-target-feature" "-xnack".* "-fsanitize=address"}}
+// XNACKNEG-NOT: {{"[^"]*clang[^"]*".* "-target-cpu" "gfx900".* "-mno-xnack".* "-fsanitize=address"}}
// XNACKNEG-NOT: {{"[^"]*clang[^"]*".* "-target-cpu" "gfx906".* "-fsanitize=address"}}
-// XNACKNEG-NOT: {{"[^"]*clang[^"]*".* "-mlink-bitcode-file" ".*asanrtl.bc".* "-target-cpu" "gfx900".* "-target-feature" "-xnack"}}
+// XNACKNEG-NOT: {{"[^"]*clang[^"]*".* "-mlink-bitcode-file" ".*asanrtl.bc".* "-target-cpu" "gfx900".* "-mno-xnack"}}
// XNACKNEG-NOT: {{"[^"]*clang[^"]*".* "-mlink-bitcode-file" ".*asanrtl.bc".* "-target-cpu" "gfx906"}}
// XNACKNEG-NOT: {{"[^"]*lld(\.exe){0,1}".* ".*hip.bc"}}
-// NOGPU-DAG: {{"[^"]*clang[^"]*".* "-target-cpu" "gfx900".* "-target-feature" "\+xnack"}}
-// NOGPU-DAG: {{"[^"]*clang[^"]*".* "-target-cpu" "gfx900".* "-target-feature" "-xnack"}}
+// NOGPU-DAG: {{"[^"]*clang[^"]*".* "-target-cpu" "gfx900".* "-mxnack"}}
+// NOGPU-DAG: {{"[^"]*clang[^"]*".* "-target-cpu" "gfx900".* "-mno-xnack"}}
// NOGPU-DAG: {{"[^"]*clang[^"]*".* "-target-cpu" "gfx906"}}
// NOGPU-DAG: {{"[^"]*clang[^"]*".* "-triple" "x86_64-unknown-linux-gnu".* "-fsanitize=address,leak"}}
// NOGPUNEG-NOT: warning: ignoring '-fsanitize=leak' option as it is not currently supported for target 'amdgcn-amd-amdhsa'
// NOGPUNEG-NOT: warning: ignoring '-fsanitize=address' option for offload arch 'gfx900:xnack-' as it is not currently supported there. Use it with an offload arch containing 'xnack+' instead
// NOGPUNEG-NOT: warning: ignoring '-fsanitize=address' option for offload arch 'gfx906' as it is not currently supported there. Use it with an offload arch containing 'xnack+' instead
-// NOGPUNEG-NOT: {{"[^"]*clang[^"]*".* "-mlink-bitcode-file" ".*asanrtl.bc".* "-target-cpu" "gfx900".* "-target-feature" "\+xnack".* "-fsanitize=address"}}
-// NOGPUNEG-NOT: {{"[^"]*clang[^"]*".* "-target-cpu" "gfx900".* "-target-feature" "\+xnack".* "-fsanitize=address,leak"}}
-// NOGPUNEG-NOT: {{"[^"]*clang[^"]*".* "-target-cpu" "gfx900".* "-target-feature" "-xnack".* "-fsanitize=address,leak"}}
+// NOGPUNEG-NOT: {{"[^"]*clang[^"]*".* "-mlink-bitcode-file" ".*asanrtl.bc".* "-target-cpu" "gfx900".* "-mxnack".* "-fsanitize=address"}}
+// NOGPUNEG-NOT: {{"[^"]*clang[^"]*".* "-target-cpu" "gfx900".* "-mxnack".* "-fsanitize=address,leak"}}
+// NOGPUNEG-NOT: {{"[^"]*clang[^"]*".* "-target-cpu" "gfx900".* "-mno-xnack".* "-fsanitize=address,leak"}}
// NOGPUNEG-NOT: {{"[^"]*clang[^"]*".* "-target-cpu" "gfx906".* "-fsanitize=address,leak"}}
-// NOGPUNEG-NOT: {{"[^"]*clang[^"]*".* "-target-cpu" "gfx900".* "-target-feature" "\+xnack".* "-fsanitize=address"}}
-// NOGPUNEG-NOT: {{"[^"]*clang[^"]*".* "-target-cpu" "gfx900".* "-target-feature" "-xnack".* "-fsanitize=address"}}
+// NOGPUNEG-NOT: {{"[^"]*clang[^"]*".* "-target-cpu" "gfx900".* "-mxnack".* "-fsanitize=address"}}
+// NOGPUNEG-NOT: {{"[^"]*clang[^"]*".* "-target-cpu" "gfx900".* "-mno-xnack".* "-fsanitize=address"}}
// NOGPUNEG-NOT: {{"[^"]*clang[^"]*".* "-target-cpu" "gfx906".* "-fsanitize=address"}}
-// NOGPUNEG-NOT: {{"[^"]*clang[^"]*".* "-mlink-bitcode-file" ".*asanrtl.bc".* "-target-cpu" "gfx900".* "-target-feature" "\+xnack"}}
-// NOGPUNEG-NOT: {{"[^"]*clang[^"]*".* "-mlink-bitcode-file" ".*asanrtl.bc".* "-target-cpu" "gfx900".* "-target-feature" "-xnack"}}
+// NOGPUNEG-NOT: {{"[^"]*clang[^"]*".* "-mlink-bitcode-file" ".*asanrtl.bc".* "-target-cpu" "gfx900".* "-mxnack"}}
+// NOGPUNEG-NOT: {{"[^"]*clang[^"]*".* "-mlink-bitcode-file" ".*asanrtl.bc".* "-target-cpu" "gfx900".* "-mno-xnack"}}
// NOGPUNEG-NOT: {{"[^"]*clang[^"]*".* "-mlink-bitcode-file" ".*asanrtl.bc".* "-target-cpu" "gfx906"}}
// NOGPUNEG-NOT: {{"[^"]*lld(\.exe){0,1}".* ".*hip.bc"}}
// INVALIDCOMBINATION1-DAG: warning: ignoring 'fuzzer' in '-fsanitize=address,fuzzer' option as it is not currently supported for target 'amdgcn-amd-amdhsa' [-Woption-ignored]
// INVALIDCOMBINATION2-DAG: warning: ignoring 'fuzzer' in '-fsanitize=fuzzer,address' option as it is not currently supported for target 'amdgcn-amd-amdhsa' [-Woption-ignored]
-// INVALIDCOMBINATION-DAG: {{"[^"]*clang[^"]*".* "-mlink-bitcode-file" ".*asanrtl.bc".* "-target-cpu" "gfx900".* "-target-feature" "\+xnack".* "-fsanitize=address"}}
+// INVALIDCOMBINATION-DAG: {{"[^"]*clang[^"]*".* "-mlink-bitcode-file" ".*asanrtl.bc".* "-target-cpu" "gfx900".* "-mxnack".* "-fsanitize=address"}}
// INVALIDCOMBINATION-DAG: {{"[^"]*clang[^"]*".* "-triple" "x86_64-unknown-linux-gnu".* "-fsanitize=address,fuzzer,fuzzer-no-link"}}
// MULT1: warning: ignoring 'leak' in '-fsanitize=leak' option as it is not currently supported for target 'amdgcn-amd-amdhsa' [-Woption-ignored]
@@ -167,7 +167,7 @@
// FIXME: This should produce a separate warning for address and fuzzer. The xnack+ hint only applies to the address part
// MULT2: warning: ignoring '-fsanitize=fuzzer,address' option for offload arch 'gfx908:xnack-' as it is not currently supported there. Use it with an offload arch containing 'xnack+' instead [-Woption-ignored]
-// XNACK2-DAG: {{"[^"]*clang[^"]*".* "-mlink-bitcode-file" ".*asanrtl.bc".* "-target-cpu" "gfx900".* "-target-feature" "\+xnack".* "-fsanitize=address"}}
+// XNACK2-DAG: {{"[^"]*clang[^"]*".* "-mlink-bitcode-file" ".*asanrtl.bc".* "-target-cpu" "gfx900".* "-mxnack".* "-fsanitize=address"}}
// XNACK2-DAG: {{"[^"]*clang[^"]*".* "-target-cpu" "gfx908"}}
// XNACK2-DAG: {{"[^"]*clang[^"]*".* "-triple" "x86_64-unknown-linux-gnu".* "-fsanitize=address,fuzzer,fuzzer-no-link,leak"}}
diff --git a/clang/test/Driver/hip-target-id.hip b/clang/test/Driver/hip-target-id.hip
index 1a1363d577d27..b465f9cfb3386 100644
--- a/clang/test/Driver/hip-target-id.hip
+++ b/clang/test/Driver/hip-target-id.hip
@@ -23,24 +23,22 @@
// CHECK: [[CLANG:"[^"]*clang[^"]*"]] "-cc1" "-triple" "amdgcn-amd-amdhsa"
// CHECK-SAME: "-target-cpu" "gfx908"
-// CHECK-SAME: "-target-feature" "+xnack"
-// CHECK-SAME: "-target-feature" "+sramecc"
+// CHECK-SAME: "-mxnack"
+// CHECK-SAME: "-msramecc"
// TMP: [[CLANG:"[^"]*clang[^"]*"]] "-cc1as" "-triple" "amdgcn-amd-amdhsa"
// TMP-SAME: "-target-cpu" "gfx908"
-// TMP-SAME: "-target-feature" "+xnack"
-// TMP-SAME: "-target-feature" "+sramecc"
+// TMP-SAME: "-mxnack"
+// TMP-SAME: "-msramecc"
// CHECK: [[LLD:"[^"]*lld[^"]*"]] {{.*}} "-plugin-opt=mcpu=gfx908"
-// CHECK-SAME: "-plugin-opt=-mattr=+xnack,+sramecc"
// CHECK: [[CLANG]] "-cc1" "-triple" "amdgcn-amd-amdhsa"
// CHECK-SAME: "-target-cpu" "gfx908"
-// CHECK-SAME: "-target-feature" "+xnack"
-// CHECK-SAME: "-target-feature" "-sramecc"
+// CHECK-SAME: "-mxnack"
+// CHECK-SAME: "-mno-sramecc"
// CHECK: [[LLD]] {{.*}} "-plugin-opt=mcpu=gfx908"
-// CHECK-SAME: "-plugin-opt=-mattr=+xnack,-sramecc"
// CHECK: {{"[^"]*clang-offload-bundler[^"]*"}}
// CHECK-SAME: "-targets=host-x86_64-unknown-linux-gnu,hipv4-amdgcn-amd-amdhsa--gfx908:sramecc+:xnack+,hipv4-amdgcn-amd-amdhsa--gfx908:sramecc-:xnack+"
diff --git a/clang/test/Driver/hip-toolchain-features.hip b/clang/test/Driver/hip-toolchain-features.hip
index d0ad0f2af4c3d..d15367723d19a 100644
--- a/clang/test/Driver/hip-toolchain-features.hip
+++ b/clang/test/Driver/hip-toolchain-features.hip
@@ -5,10 +5,8 @@
// RUN: -nogpuinc --offload-arch=gfx906:xnack- --offload-arch=gfx900:xnack- %s \
// RUN: 2>&1 | FileCheck %s -check-prefix=NOXNACK
-// XNACK: {{.*}}clang{{.*}}"-target-feature" "+xnack"
-// NOXNACK: {{.*}}clang{{.*}}"-target-feature" "-xnack"
-// XNACK: {{.*}}lld{{.*}} "-plugin-opt=-mattr=+xnack"
-// NOXNACK: {{.*}}lld{{.*}} "-plugin-opt=-mattr=-xnack"
+// XNACK: {{.*}}clang{{.*}}"-mxnack"
+// NOXNACK: {{.*}}clang{{.*}}"-mno-xnack"
// RUN: %clang -### --target=x86_64-linux-gnu -fgpu-rdc -nogpulib \
// RUN: -nogpuinc --offload-arch=gfx908:sramecc+ --no-offload-new-driver %s \
@@ -17,10 +15,8 @@
// RUN: -nogpuinc --offload-arch=gfx908:sramecc- --no-offload-new-driver %s \
// RUN: 2>&1 | FileCheck %s -check-prefix=NOSRAM
-// SRAM: {{.*}}clang{{.*}}"-target-feature" "+sramecc"
-// NOSRAM: {{.*}}clang{{.*}}"-target-feature" "-sramecc"
-// SRAM: {{.*}}lld{{.*}} "-plugin-opt=-mattr=+sramecc"
-// NOTSRAM: {{.*}}lld{{.*}} "-plugin-opt=-mattr=-sramecc"
+// SRAM: {{.*}}clang{{.*}}"-msramecc"
+// NOSRAM: {{.*}}clang{{.*}}"-mno-sramecc"
// RUN: %clang -### --target=x86_64-linux-gnu -fgpu-rdc -nogpulib \
// RUN: -nogpuinc --offload-arch=gfx1010 --no-offload-new-driver %s \
@@ -41,8 +37,8 @@
// RUN: -nogpuinc --offload-arch=gfx908:xnack-:sramecc- --no-offload-new-driver %s \
// RUN: 2>&1 | FileCheck %s -check-prefix=NOALL3
-// ALL3: {{.*}}clang{{.*}}"-target-feature" "+xnack" "-target-feature" "+sramecc"
-// NOALL3: {{.*}}clang{{.*}}"-target-feature" "-xnack" "-target-feature" "-sramecc"
+// ALL3: {{.*}}clang{{.*}}"-mxnack" {{.*}}"-msramecc"
+// NOALL3: {{.*}}clang{{.*}}"-mno-xnack" {{.*}}"-mno-sramecc"
// RUN: %clang -### --target=x86_64-linux-gnu -fgpu-rdc -nogpulib \
// RUN: -nogpuinc --offload-arch=gfx1010 --no-offload-new-driver %s \
diff --git a/clang/test/Driver/target-id.cl b/clang/test/Driver/target-id.cl
index 685d5f8665b63..45a57bdc90abf 100644
--- a/clang/test/Driver/target-id.cl
+++ b/clang/test/Driver/target-id.cl
@@ -8,7 +8,7 @@
// RUN: %clang -### -target amdgcn-amd-amdhsa \
// RUN: -mcpu=gfx908:xnack+:sramecc- \
-// RUN: -nostdlib -x assembler %s 2>&1 | FileCheck %s
+// RUN: -nostdlib -x assembler %s 2>&1 | FileCheck -check-prefix=ASM %s
// RUN: %clang -### -target amdgcn-amd-amdpal \
// RUN: -mcpu=gfx908:xnack+:sramecc- \
@@ -22,8 +22,12 @@
// RUN: -nostdlib %s 2>&1 | FileCheck -check-prefix=NONE %s
// CHECK: "-target-cpu" "gfx908"
-// CHECK-SAME: "-target-feature" "+xnack"
-// CHECK-SAME: "-target-feature" "-sramecc"
+// CHECK-SAME: "-mxnack"
+// CHECK-SAME: "-mno-sramecc"
+
+// ASM: "-target-cpu" "gfx908"
+// ASM-NOT: "-mxnack"
+// ASM-NOT: "-mno-sramecc"
// NONE-NOT: "-target-cpu"
// NONE-NOT: "-target-feature"
diff --git a/llvm/docs/AMDGPUUsage.rst b/llvm/docs/AMDGPUUsage.rst
index dba0997e4f099..aeb7d0436091c 100644
--- a/llvm/docs/AMDGPUUsage.rst
+++ b/llvm/docs/AMDGPUUsage.rst
@@ -871,10 +871,48 @@ consumed by the AMDGPU backend during code generation.
- Same as above, but for typed buffer instructions (``tbuffer_load`` /
``tbuffer_store``).
+ * - ``amdgpu.xnack``
+ - ``i32``
+ - Error
+ - Controls XNACK (page fault) replay mode. This is ignored on
+ targets which do not support xnack.
+
+ - absent: **any**. The module can be loaded and executed in a process
+ with XNACK replay either enabled or disabled. Code generation
+ assumes XNACK may be enabled.
+ - ``0``: **off**. The module can only be loaded and executed in a
+ process with XNACK replay disabled. Code generation is optimized
+ for XNACK disabled.
+ - ``1``: **on**. The module can only be loaded and executed in a
+ process with XNACK replay enabled. Code generation assumes XNACK
+ is enabled.
+
+ At link time, modules with conflicting settings (``0`` vs ``1``)
+ produce an error. Modules with **any** (absent flag) are compatible
+ with any setting.
+
+ * - ``amdgpu.sramecc``
+ - ``i32``
+ - Error
+ - Controls SRAMECC mode. This is ignored on targets which do not
+ support sramecc.
+
+ - absent: **any**. The module can be loaded and executed in a process
+ with SRAMECC either enabled or disabled.
+ - ``0``: **off**. The module can only be loaded and executed in a
+ process with SRAMECC disabled.
+ - ``1``: **on**. The module can only be loaded and executed in a
+ process with SRAMECC enabled. Some instructions behave differently
+ (e.g., D16 memory instructions).
+
+ At link time, modules with conflicting settings (``0`` vs ``1``)
+ produce an error. Modules with **any** (absent flag) are compatible
+ with any setting.
+
.. note::
Frontends that require misaligned-access merging for performance should
- set both flags to ``1`` (relaxed). Frontends that require strict
+ set both buffer OOB flags to ``1`` (relaxed). Frontends that require strict
per-byte OOB guarantees should set the flags to ``2`` (strict) as needed.
Modules that do not use buffer operations or are indifferent to OOB semantics
(e.g. device libraries) should leave the flags absent.
diff --git a/llvm/docs/ReleaseNotes.md b/llvm/docs/ReleaseNotes.md
index df4ced91e8f5e..061564e4e646d 100644
--- a/llvm/docs/ReleaseNotes.md
+++ b/llvm/docs/ReleaseNotes.md
@@ -197,6 +197,8 @@ Makes programs 10x faster by doing Special New Thing.
previously relied on the subtarget feature to enable misaligned buffer merging
must now set the corresponding module flag to `1` (relaxed). An absent flag is
treated as strict by the backend.
+* Replaced `xnack` and `sramecc` target features with `amdgpu.xnack`
+ and `amdgpu.sramecc` module flags.
### Changes to the ARM Backend
diff --git a/llvm/lib/IR/Verifier.cpp b/llvm/lib/IR/Verifier.cpp
index 14b1f7d799ea9..c38870346794b 100644
--- a/llvm/lib/IR/Verifier.cpp
+++ b/llvm/lib/IR/Verifier.cpp
@@ -2160,6 +2160,25 @@ Verifier::visitModuleFlag(const MDNode *Op,
Check(Value->getZExtValue() <= 2,
"'" + ID->getString() + "' module flag must be 0, 1, or 2");
}
+
+ return;
+ }
+
+ if (ID->getString() == "amdgpu.xnack" ||
+ ID->getString() == "amdgpu.sramecc") {
+ Check(MFB == Module::Error,
+ "'" + ID->getString() +
+ "' module flag must use 'error' merge behaviour");
+ ConstantInt *Value =
+ mdconst::dyn_extract_or_null<ConstantInt>(Op->getOperand(2));
+ Check(Value, "'" + ID->getString() +
+ "' module flag must have a constant integer value");
+ if (Value) {
+ Check(Value->getZExtValue() <= 1,
+ "'" + ID->getString() + "' module flag must be 0 or 1");
+ }
+
+ return;
}
if (ID->getString() == "CG Profile") {
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp b/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
index 9f29934f5a4a5..2510b30ece201 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
@@ -194,30 +194,6 @@ void AMDGPUAsmPrinter::emitFunctionBodyStart() {
if (!getTargetStreamer()->getTargetID())
initializeTargetID(*F.getParent());
- const auto &FunctionTargetID = STM.getTargetID();
- // Make sure function's xnack settings are compatible with module's
- // xnack settings.
- if (FunctionTargetID.isXnackSupported() &&
- FunctionTargetID.getXnackSetting() != IsaInfo::TargetIDSetting::Any &&
- FunctionTargetID.getXnackSetting() !=
- getTargetStreamer()->getTargetID()->getXnackSetting()) {
- OutContext.reportError(
- {}, "xnack setting of '" + Twine(MF->getName()) +
- "' function does not match module xnack setting");
- return;
- }
- // Make sure function's sramecc settings are compatible with module's
- // sramecc settings.
- if (FunctionTargetID.isSramEccSupported() &&
- FunctionTargetID.getSramEccSetting() != IsaInfo::TargetIDSetting::Any &&
- FunctionTargetID.getSramEccSetting() !=
- getTargetStreamer()->getTargetID()->getSramEccSetting()) {
- OutContext.reportError(
- {}, "sramecc setting of '" + Twine(MF->getName()) +
- "' function does not match module sramecc setting");
- return;
- }
-
if (!MFI.isEntryFunction())
return;
@@ -1194,31 +1170,39 @@ void AMDGPUAsmPrinter::emitDVgprSymbol(MachineFunction &MF) {
// TODO: Fold this into emitFunctionBodyStart.
void AMDGPUAsmPrinter::initializeTargetID(const Module &M) {
- // In the beginning all features are either 'Any' or 'NotSupported',
- // depending on global target features. This will cover empty modules.
- getTargetStreamer()->initializeTargetID(*getGlobalSTI(),
- getGlobalSTI()->getFeatureString());
-
- // If module is empty, we are done.
- if (M.empty())
- return;
-
- // If module is not empty, need to find first 'Off' or 'On' feature
- // setting per feature from functions in module.
- for (auto &F : M) {
- auto &TSTargetID = getTargetStreamer()->getTargetID();
- if ((!TSTargetID->isXnackSupported() || TSTargetID->isXnackOnOrOff()) &&
- (!TSTargetID->isSramEccSupported() || TSTargetID->isSramEccOnOrOff()))
- break;
-
- const GCNSubtarget &STM = TM.getSubtarget<GCNSubtarget>(F);
- const IsaInfo::AMDGPUTargetID &STMTargetID = STM.getTargetID();
- if (TSTargetID->isXnackSupported())
- if (TSTargetID->getXnackSetting() == IsaInfo::TargetIDSetting::Any)
- TSTargetID->setXnackSetting(STMTargetID.getXnackSetting());
- if (TSTargetID->isSramEccSupported())
- if (TSTargetID->getSramEccSetting() == IsaInfo::TargetIDSetting::Any)
- TSTargetID->setSramEccSetting(STMTargetID.getSramEccSetting());
+ getTargetStreamer()->initializeTargetID(*getGlobalSTI());
+
+ auto &TSTargetID = getTargetStreamer()->getTargetID();
+
+ // Error if -mattr specified xnack or sramecc.
+ // TODO: Remove this when subtarget features removed.
+ StringRef FeatureString = getGlobalSTI()->getFeatureString();
+ if (FeatureString.contains("xnack")) {
+ M.getContext().diagnose(DiagnosticInfoGeneric(
+ "xnack/sramecc should be specified via module flags. "
+ "Use module flag 'amdgpu.xnack' instead of subtarget feature",
+ DS_Error));
+ }
+ if (FeatureString.contains("sramecc")) {
+ M.getContext().diagnose(DiagnosticInfoGeneric(
+ "xnack/sramecc should be specified via module flags. "
+ "Use module flag 'amdgpu.sramecc' instead of subtarget feature",
+ DS_Error));
+ }
+
+ // Apply xnack/sramecc settings from module flags.
+ if (getGlobalSTI()->getFeatureBits().test(AMDGPU::FeatureXNACKOnOffModes)) {
+ AMDGPU::IsaInfo::TargetIDSetting Setting =
+ GCNTargetMachine::getTargetIDSettingFromModuleFlag(M, "amdgpu.xnack");
+ if (Setting != IsaInfo::TargetIDSetting::Any)
+ TSTargetID->setXnackSetting(Setting);
+ }
+
+ if (getGlobalSTI()->getFeatureBits().test(AMDGPU::FeatureSupportsSRAMECC)) {
+ AMDGPU::IsaInfo::TargetIDSetting Setting =
+ GCNTargetMachine::getTargetIDSettingFromModuleFlag(M, "amdgpu.sramecc");
+ if (Setting != IsaInfo::TargetIDSetting::Any)
+ TSTargetID->setSramEccSetting(Setting);
}
}
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp b/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
index b078e0835a90e..9fa371602e82d 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
@@ -468,6 +468,16 @@ static cl::opt<bool> EnableAMDGPUAliasAnalysis("enable-amdgpu-aa", cl::Hidden,
cl::desc("Enable AMDGPU Alias Analysis"),
cl::init(true));
+static cl::opt<bool>
+ XnackSetting("amdgpu-xnack",
+ cl::desc("Force amdgpu.xnack value for testing"),
+ cl::ReallyHidden);
+
+static cl::opt<bool>
+ SramEccSetting("amdgpu-sramecc",
+ cl::desc("Force amdgpu.sramecc for testing"),
+ cl::ReallyHidden);
+
// Enable lib calls simplifications
static cl::opt<bool> EnableLibCallSimplify(
"amdgpu-simplify-libcall",
@@ -1262,6 +1272,26 @@ static OOBFlagValue getOOBFlagValue(const Module &M, StringRef FlagName) {
return static_cast<OOBFlagValue>(Flag->getZExtValue());
}
+/// Returns the xnack/sramecc setting encoded by a module flag.
+/// Module flag values: 0 = disabled, 1 = enabled.
+/// An absent flag defaults to Any.
+AMDGPU::IsaInfo::TargetIDSetting
+GCNTargetMachine::getTargetIDSettingFromModuleFlag(const Module &M,
+ StringRef FlagName) {
+ using AMDGPU::IsaInfo::TargetIDSetting;
+
+ if (XnackSetting.getNumOccurrences() > 0 && FlagName == "amdgpu.xnack")
+ return XnackSetting ? TargetIDSetting::On : TargetIDSetting::Off;
+ if (SramEccSetting.getNumOccurrences() > 0 && FlagName == "amdgpu.sramecc")
+ return SramEccSetting ? TargetIDSetting::On : TargetIDSetting::Off;
+
+ const auto *Flag =
+ mdconst::dyn_extract_or_null<ConstantInt>(M.getModuleFlag(FlagName));
+ if (!Flag)
+ return TargetIDSetting::Any;
+ return Flag->getZExtValue() == 0 ? TargetIDSetting::Off : TargetIDSetting::On;
+}
+
const TargetSubtargetInfo *
GCNTargetMachine::getSubtargetImpl(const Function &F) const {
StringRef GPU = getGPUName(F);
@@ -1272,17 +1302,31 @@ GCNTargetMachine::getSubtargetImpl(const Function &F) const {
OOBFlagValue TBufOOB = getOOBFlagValue(M, AMDGPUOOBMode::TBufferFlag);
bool BufRelaxed = BufOOB == OOBFlagValue::Relaxed;
bool TBufRelaxed = TBufOOB == OOBFlagValue::Relaxed;
+
+ using AMDGPU::IsaInfo::TargetIDSetting;
+ TargetIDSetting Xnack = getTargetIDSettingFromModuleFlag(M, "amdgpu.xnack");
+ TargetIDSetting SramEcc =
+ getTargetIDSettingFromModuleFlag(M, "amdgpu.sramecc");
+
SmallString<128> SubtargetKey(GPU);
SubtargetKey.append(FS);
if (BufRelaxed)
SubtargetKey.append(",buf-oob=1");
if (TBufRelaxed)
SubtargetKey.append(",tbuf-oob=1");
+ if (Xnack != TargetIDSetting::Any) {
+ SubtargetKey.append(",xnack=");
+ SubtargetKey.push_back(Xnack == TargetIDSetting::On ? '1' : '0');
+ }
+ if (SramEcc != TargetIDSetting::Any) {
+ SubtargetKey.append(",sramecc=");
+ SubtargetKey.push_back(Xnack == TargetIDSetting::On ? '1' : '0');
+ }
auto &I = SubtargetMap[SubtargetKey];
if (!I) {
I = std::make_unique<GCNSubtarget>(TargetTriple, GPU, FS, *this, BufRelaxed,
- TBufRelaxed);
+ TBufRelaxed, Xnack, SramEcc);
}
I->setScalarizeGlobalBehavior(ScalarizeGlobal);
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.h b/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.h
index e2c27f3822380..25acbd747201d 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.h
@@ -107,6 +107,11 @@ class GCNTargetMachine final : public AMDGPUTargetMachine {
void registerMachineRegisterInfoCallback(MachineFunction &MF) const override;
+ /// Get xnack/sramecc setting from module flag or cl::opt (for testing).
+ /// Returns Any if not specified.
+ static AMDGPU::IsaInfo::TargetIDSetting
+ getTargetIDSettingFromModuleFlag(const Module &M, StringRef FlagName);
+
MachineFunctionInfo *
createMachineFunctionInfo(BumpPtrAllocator &Allocator, const Function &F,
const TargetSubtargetInfo *STI) const override;
diff --git a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
index b1ef7d93e4363..41d0eae0adc71 100644
--- a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
+++ b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
@@ -9402,7 +9402,7 @@ void AMDGPUAsmParser::onBeginOfFile() {
if (!getTargetStreamer().getTargetID())
getTargetStreamer().initializeTargetID(getSTI(),
- getSTI().getFeatureString());
+ /*ApplyFeatureString=*/true);
if (isHsaAbi(getSTI()))
getTargetStreamer().EmitDirectiveAMDGCNTarget();
diff --git a/llvm/lib/Target/AMDGPU/GCNSubtarget.cpp b/llvm/lib/Target/AMDGPU/GCNSubtarget.cpp
index 55edfc2ea52d2..b90e0812372a3 100644
--- a/llvm/lib/Target/AMDGPU/GCNSubtarget.cpp
+++ b/llvm/lib/Target/AMDGPU/GCNSubtarget.cpp
@@ -157,11 +157,6 @@ GCNSubtarget &GCNSubtarget::initializeSubtargetDependencies(const Triple &TT,
assert(llvm::isPowerOf2_32(InstCacheLineSize) &&
"InstCacheLineSize must be a power of 2");
- LLVM_DEBUG(dbgs() << "xnack setting for subtarget: "
- << TargetID.getXnackSetting() << '\n');
- LLVM_DEBUG(dbgs() << "sramecc setting for subtarget: "
- << TargetID.getSramEccSetting() << '\n');
-
return *this;
}
@@ -176,11 +171,13 @@ void GCNSubtarget::checkSubtargetFeatures(const Function &F) const {
GCNSubtarget::GCNSubtarget(const Triple &TT, StringRef GPU, StringRef FS,
const GCNTargetMachine &TM, bool BufferOOBRelaxed,
- bool TBufferOOBRelaxed)
+ bool TBufferOOBRelaxed,
+ AMDGPU::IsaInfo::TargetIDSetting XnackSetting,
+ AMDGPU::IsaInfo::TargetIDSetting SramEccSetting)
: // clang-format off
AMDGPUGenSubtargetInfo(TT, GPU, /*TuneCPU*/ GPU, FS),
AMDGPUSubtarget(TT),
- TargetID(*this, FS),
+ TargetID(*this),
InstrItins(getInstrItineraryForCPU(GPU)),
BufferOOBRelaxed(BufferOOBRelaxed),
TBufferOOBRelaxed(TBufferOOBRelaxed),
@@ -190,6 +187,22 @@ GCNSubtarget::GCNSubtarget(const Triple &TT, StringRef GPU, StringRef FS,
FrameLowering(TargetFrameLowering::StackGrowsUp, getStackAlignment(), 0,
/*TransAl=*/Align(4)) {
// clang-format on
+
+ // Apply the module flag's xnack setting if the target supports on/off modes.
+ // Targets without on/off mode support have xnack always on and ignore module
+ // flags.
+ if (hasXNACKOnOffModes())
+ TargetID.setXnackSetting(XnackSetting);
+
+ // Apply the module flag's sramecc setting if the target supports it.
+ if (supportsSRAMECC())
+ TargetID.setSramEccSetting(SramEccSetting);
+
+ LLVM_DEBUG(dbgs() << "xnack setting for subtarget: "
+ << TargetID.getXnackSetting() << '\n');
+ LLVM_DEBUG(dbgs() << "sramecc setting for subtarget: "
+ << TargetID.getSramEccSetting() << '\n');
+
MaxWavesPerEU = AMDGPU::IsaInfo::getMaxWavesPerEU(*this);
EUsPerCU = AMDGPU::IsaInfo::getEUsPerCU(*this);
diff --git a/llvm/lib/Target/AMDGPU/GCNSubtarget.h b/llvm/lib/Target/AMDGPU/GCNSubtarget.h
index af47a8725c2d0..f30ca906aaf6e 100644
--- a/llvm/lib/Target/AMDGPU/GCNSubtarget.h
+++ b/llvm/lib/Target/AMDGPU/GCNSubtarget.h
@@ -113,7 +113,11 @@ class GCNSubtarget final : public AMDGPUGenSubtargetInfo,
public:
GCNSubtarget(const Triple &TT, StringRef GPU, StringRef FS,
const GCNTargetMachine &TM, bool BufferOOBRelaxed = false,
- bool TBufferOOBRelaxed = false);
+ bool TBufferOOBRelaxed = false,
+ AMDGPU::IsaInfo::TargetIDSetting XnackSetting =
+ AMDGPU::IsaInfo::TargetIDSetting::Any,
+ AMDGPU::IsaInfo::TargetIDSetting SramEccSetting =
+ AMDGPU::IsaInfo::TargetIDSetting::Any);
~GCNSubtarget() override;
GCNSubtarget &initializeSubtargetDependencies(const Triple &TT, StringRef GPU,
diff --git a/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.cpp b/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.cpp
index e9d7cc3f1476d..35f83d5026fb9 100644
--- a/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.cpp
+++ b/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.cpp
@@ -46,6 +46,14 @@ static cl::opt<unsigned>
"added. For testing purposes only."),
cl::ReallyHidden, cl::init(0));
+void AMDGPUTargetStreamer::initializeTargetID(const MCSubtargetInfo &STI,
+ bool ApplyFeatureString) {
+ assert(TargetID == std::nullopt && "TargetID can only be initialized once");
+ // Apply xnack/sramecc from subtarget features only in MC contexts
+ // (assembler), not in codegen where they come from module flags
+ TargetID.emplace(STI, ApplyFeatureString ? STI.getFeatureString() : "");
+}
+
bool AMDGPUTargetStreamer::EmitHSAMetadataV3(StringRef HSAMetadataString) {
msgpack::Document HSAMetadataDoc;
if (!HSAMetadataDoc.fromYAML(HSAMetadataString))
diff --git a/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.h b/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.h
index dc9636c6c2105..d41bf3c4c6712 100644
--- a/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.h
+++ b/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.h
@@ -139,10 +139,8 @@ class AMDGPUTargetStreamer : public MCTargetStreamer {
std::optional<AMDGPU::IsaInfo::AMDGPUTargetID> &getTargetID() {
return TargetID;
}
- void initializeTargetID(const MCSubtargetInfo &STI, StringRef FeatureString) {
- assert(TargetID == std::nullopt && "TargetID can only be initialized once");
- TargetID.emplace(STI, FeatureString);
- }
+ void initializeTargetID(const MCSubtargetInfo &STI,
+ bool ApplyFeatureString = false);
};
class AMDGPUTargetAsmStreamer final : public AMDGPUTargetStreamer {
diff --git a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
index e1e83ece32ad0..8428b0aeeb625 100644
--- a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
@@ -1099,15 +1099,17 @@ VOPD::InstInfo getVOPDInstInfo(unsigned VOPDOpcode,
namespace IsaInfo {
-AMDGPUTargetID::AMDGPUTargetID(const MCSubtargetInfo &STI,
- StringRef FeatureString)
+AMDGPUTargetID::AMDGPUTargetID(const MCSubtargetInfo &STI)
: STI(STI), XnackSetting(STI.getFeatureBits().test(FeatureSupportsXNACK)
? TargetIDSetting::Any
: TargetIDSetting::Unsupported),
SramEccSetting(STI.getFeatureBits().test(FeatureSupportsSRAMECC)
? TargetIDSetting::Any
- : TargetIDSetting::Unsupported) {
+ : TargetIDSetting::Unsupported) {}
+AMDGPUTargetID::AMDGPUTargetID(const MCSubtargetInfo &STI,
+ StringRef FeatureString)
+ : AMDGPUTargetID(STI) {
// Check if xnack or sramecc is explicitly enabled or disabled. In the
// absence of the target features we assume we must generate code that can run
// in any environment.
@@ -1130,7 +1132,7 @@ AMDGPUTargetID::AMDGPUTargetID(const MCSubtargetInfo &STI,
// Targets without on/off mode support keep their initial setting (Any).
bool XnackSupported = STI.getFeatureBits().test(FeatureXNACKOnOffModes);
- bool SramEccSupported = isSramEccSupported();
+ bool SramEccSupported = STI.getFeatureBits().test(FeatureSupportsSRAMECC);
if (XnackRequested) {
if (XnackSupported) {
@@ -1178,6 +1180,26 @@ getTargetIDSettingFromFeatureString(StringRef FeatureString) {
llvm_unreachable("Malformed feature string");
}
+std::optional<std::pair<StringRef, StringRef>>
+AMDGPUTargetID::parseTargetIDDirective(StringRef TargetIDDirective) {
+ // Extract the CPU+features from the target ID directive
+ // Format: "amdgcn-amd-amdhsa--gfx1010:xnack-:sramecc+"
+ // Split on '-' and take the last element
+ SmallVector<StringRef, 8> Parts;
+ TargetIDDirective.split(Parts, '-');
+
+ if (Parts.empty())
+ return std::nullopt;
+
+ // The CPU+features is the last element after splitting on '-'
+ StringRef CPUAndFeatures = Parts.back();
+
+ // Extract CPU name (split on ':' and take the first part)
+ StringRef CPUName = CPUAndFeatures.split(':').first;
+
+ return std::make_pair(CPUName, CPUAndFeatures);
+}
+
void AMDGPUTargetID::setTargetIDFromTargetIDStream(StringRef TargetID) {
SmallVector<StringRef, 3> TargetIDSplit;
TargetID.split(TargetIDSplit, ':');
@@ -1190,6 +1212,27 @@ void AMDGPUTargetID::setTargetIDFromTargetIDStream(StringRef TargetID) {
}
}
+// FIXME: This is a hack which should be removed. This attempts to turn these
+// feature bits into tri-states, where unspecified is distinct from absent. The
+// feature string is supposed to be evaluated as a bitvector only.
+void AMDGPUTargetID::setTargetIDFromFeaturesString(StringRef Features) {
+ // Parse features like "+xnack", "-xnack", "+sramecc", "-sramecc"
+ SmallVector<StringRef, 16> FeatureList;
+ Features.split(FeatureList, ',');
+
+ for (StringRef Feature : FeatureList) {
+ Feature = Feature.trim();
+ if (Feature == "+xnack")
+ XnackSetting = TargetIDSetting::On;
+ else if (Feature == "-xnack")
+ XnackSetting = TargetIDSetting::Off;
+ else if (Feature == "+sramecc")
+ SramEccSetting = TargetIDSetting::On;
+ else if (Feature == "-sramecc")
+ SramEccSetting = TargetIDSetting::Off;
+ }
+}
+
void AMDGPUTargetID::print(raw_ostream &StreamRep) const {
const Triple &TargetTriple = STI.getTargetTriple();
auto Version = getIsaVersion(STI.getCPU());
diff --git a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h
index 6c771b3460662..5aa44f8de4b20 100644
--- a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h
+++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h
@@ -160,6 +160,7 @@ class AMDGPUTargetID {
TargetIDSetting SramEccSetting;
public:
+ explicit AMDGPUTargetID(const MCSubtargetInfo &STI);
explicit AMDGPUTargetID(const MCSubtargetInfo &STI, StringRef FeatureString);
~AMDGPUTargetID() = default;
@@ -219,6 +220,18 @@ class AMDGPUTargetID {
void setTargetIDFromTargetIDStream(StringRef TargetID);
+ /// Parse subtarget feature string and set xnack/sramecc settings.
+ /// Looks for "+xnack"/"-xnack" and "+sramecc"/"-sramecc" in the feature
+ /// string.
+ void setTargetIDFromFeaturesString(StringRef Features);
+
+ /// Parse a target ID directive string (e.g.,
+ /// "amdgcn-amd-amdhsa--gfx1010:xnack-") and extract the CPU name and
+ /// CPU+features string.
+ /// \returns A pair of (CPUName, CPUAndFeatures) or std::nullopt if malformed.
+ static std::optional<std::pair<StringRef, StringRef>>
+ parseTargetIDDirective(StringRef TargetIDDirective);
+
/// Write string representation to \p OS
void print(raw_ostream &OS) const;
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/extractelement-stack-lower.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/extractelement-stack-lower.ll
index a28827a8c4683..60a4a8c33422e 100644
--- a/llvm/test/CodeGen/AMDGPU/GlobalISel/extractelement-stack-lower.ll
+++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/extractelement-stack-lower.ll
@@ -1,5 +1,5 @@
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
-; RUN: llc -global-isel -new-reg-bank-select -mtriple=amdgcn-mesa-mesa3d -mcpu=gfx900 -mattr=-xnack < %s | FileCheck -check-prefixes=GFX9 %s
+; RUN: llc -global-isel -new-reg-bank-select -mtriple=amdgcn-mesa-mesa3d -mcpu=gfx900 --amdgpu-xnack=false < %s | FileCheck -check-prefixes=GFX9 %s
; RUN: llc -global-isel -new-reg-bank-select -mtriple=amdgcn-mesa-mesa3d -mcpu=gfx1200 < %s | FileCheck -check-prefixes=GFX12 %s
; Check lowering of some large extractelement that use the stack
diff --git a/llvm/test/CodeGen/AMDGPU/amdpal-callable.ll b/llvm/test/CodeGen/AMDGPU/amdpal-callable.ll
index cd92f8f60f6c4..01fbab49564e1 100644
--- a/llvm/test/CodeGen/AMDGPU/amdpal-callable.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdpal-callable.ll
@@ -1,6 +1,6 @@
-; RUN: llc -mtriple=amdgcn--amdpal -mattr=-xnack -mattr=+dx10-clamp-and-ieee-mode < %s | FileCheck -check-prefixes=GCN,SDAG,GFX8 -enable-var-scope %s
-; RUN: llc -mtriple=amdgcn--amdpal -mcpu=gfx900 -mattr=-xnack < %s | FileCheck -check-prefixes=GCN,SDAG,GFX9 -enable-var-scope %s
-; RUN: llc -global-isel -new-reg-bank-select -mtriple=amdgcn--amdpal -mattr=-xnack -mcpu=gfx900 < %s | FileCheck -check-prefixes=GCN,GISEL,GFX9 -enable-var-scope %s
+; RUN: llc -mtriple=amdgcn--amdpal --amdgpu-xnack=false -mattr=+dx10-clamp-and-ieee-mode < %s | FileCheck -check-prefixes=GCN,SDAG,GFX8 -enable-var-scope %s
+; RUN: llc -mtriple=amdgcn--amdpal -mcpu=gfx900 --amdgpu-xnack=false < %s | FileCheck -check-prefixes=GCN,SDAG,GFX9 -enable-var-scope %s
+; RUN: llc -global-isel -new-reg-bank-select -mtriple=amdgcn--amdpal --amdgpu-xnack=false -mcpu=gfx900 < %s | FileCheck -check-prefixes=GCN,GISEL,GFX9 -enable-var-scope %s
declare amdgpu_gfx float @extern_func(float) #0
declare amdgpu_gfx float @extern_func_many_args(<64 x float>) #0
diff --git a/llvm/test/CodeGen/AMDGPU/break-smem-soft-clauses.mir b/llvm/test/CodeGen/AMDGPU/break-smem-soft-clauses.mir
index a7962375d76c5..e87a11e9bbfdf 100644
--- a/llvm/test/CodeGen/AMDGPU/break-smem-soft-clauses.mir
+++ b/llvm/test/CodeGen/AMDGPU/break-smem-soft-clauses.mir
@@ -1,7 +1,7 @@
# RUN: llc -mtriple=amdgcn -mcpu=carrizo -verify-machineinstrs -run-pass post-RA-hazard-rec %s -o - | FileCheck -check-prefixes=GCN,XNACK %s
-# RUN: llc -mtriple=amdgcn -mcpu=fiji -mattr=-xnack -verify-machineinstrs -run-pass post-RA-hazard-rec %s -o - | FileCheck -check-prefixes=GCN %s
+# RUN: llc -mtriple=amdgcn -mcpu=fiji --amdgpu-xnack=false -verify-machineinstrs -run-pass post-RA-hazard-rec %s -o - | FileCheck -check-prefixes=GCN %s
-# RUN: llc -mtriple=amdgcn -mcpu=fiji -mattr=-xnack -passes post-RA-hazard-rec %s -o - | FileCheck -check-prefixes=GCN %s
+# RUN: llc -mtriple=amdgcn -mcpu=fiji --amdgpu-xnack=false -passes post-RA-hazard-rec %s -o - | FileCheck -check-prefixes=GCN %s
---
# Trivial clause at beginning of program
diff --git a/llvm/test/CodeGen/AMDGPU/callee-special-input-vgprs-packed.ll b/llvm/test/CodeGen/AMDGPU/callee-special-input-vgprs-packed.ll
index a7009c4d20e33..0b0b9782bf183 100644
--- a/llvm/test/CodeGen/AMDGPU/callee-special-input-vgprs-packed.ll
+++ b/llvm/test/CodeGen/AMDGPU/callee-special-input-vgprs-packed.ll
@@ -1,6 +1,6 @@
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
; RUN: opt -passes=amdgpu-attributor -mcpu=kaveri < %s | llc -mcpu=gfx90a -enable-ipra=0 | FileCheck -enable-var-scope -check-prefixes=GCN,GFX7 %s
-; RUN: opt -passes=amdgpu-attributor -mcpu=gfx90a -mattr=-xnack < %s | llc -mcpu=gfx90a -mattr=-xnack -enable-ipra=0 | FileCheck -enable-var-scope -check-prefixes=GCN,GFX90A %s
+; RUN: opt -passes=amdgpu-attributor -mcpu=gfx90a --amdgpu-xnack=false < %s | llc -mcpu=gfx90a --amdgpu-xnack=false -enable-ipra=0 | FileCheck -enable-var-scope -check-prefixes=GCN,GFX90A %s
target triple = "amdgcn-amd-amdhsa"
diff --git a/llvm/test/CodeGen/AMDGPU/cluster-flat-loads-postra.mir b/llvm/test/CodeGen/AMDGPU/cluster-flat-loads-postra.mir
index e42a1d15f0b2b..34b324601d9eb 100644
--- a/llvm/test/CodeGen/AMDGPU/cluster-flat-loads-postra.mir
+++ b/llvm/test/CodeGen/AMDGPU/cluster-flat-loads-postra.mir
@@ -1,5 +1,5 @@
-# RUN: llc -mtriple=amdgcn -mcpu=tonga -mattr=-xnack -run-pass post-RA-sched -verify-machineinstrs -o - %s | FileCheck -check-prefix=GCN %s
-# RUN: llc -mtriple=amdgcn -mcpu=tonga -mattr=-xnack -passes=post-RA-sched -o - %s | FileCheck -check-prefix=GCN %s
+# RUN: llc -mtriple=amdgcn -mcpu=tonga --amdgpu-xnack=false -run-pass post-RA-sched -verify-machineinstrs -o - %s | FileCheck -check-prefix=GCN %s
+# RUN: llc -mtriple=amdgcn -mcpu=tonga --amdgpu-xnack=false -passes=post-RA-sched -o - %s | FileCheck -check-prefix=GCN %s
# GCN: FLAT_LOAD_DWORD
# GCN-NEXT: FLAT_LOAD_DWORD
diff --git a/llvm/test/CodeGen/AMDGPU/cluster_stores.ll b/llvm/test/CodeGen/AMDGPU/cluster_stores.ll
index da5e73199a223..5c6668168aa22 100644
--- a/llvm/test/CodeGen/AMDGPU/cluster_stores.ll
+++ b/llvm/test/CodeGen/AMDGPU/cluster_stores.ll
@@ -1,5 +1,5 @@
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
-; RUN: llc -mtriple=amdgcn -mcpu=gfx900 -mattr=-xnack -debug-only=machine-scheduler < %s 2> %t | FileCheck --enable-var-scope --check-prefix=GFX9 %s
+; RUN: llc -mtriple=amdgcn -mcpu=gfx900 --amdgpu-xnack=false -debug-only=machine-scheduler < %s 2> %t | FileCheck --enable-var-scope --check-prefix=GFX9 %s
; RUN: FileCheck --enable-var-scope --check-prefix=DBG %s < %t
; RUN: llc -mtriple=amdgcn -mcpu=gfx1010 -debug-only=machine-scheduler < %s 2> %t | FileCheck --enable-var-scope --check-prefix=GFX10 %s
; RUN: FileCheck --enable-var-scope --check-prefix=DBG %s < %t
diff --git a/llvm/test/CodeGen/AMDGPU/directive-amdgcn-target.ll b/llvm/test/CodeGen/AMDGPU/directive-amdgcn-target.ll
index a7ac3bd81cd9f..80ced56fc1f62 100644
--- a/llvm/test/CodeGen/AMDGPU/directive-amdgcn-target.ll
+++ b/llvm/test/CodeGen/AMDGPU/directive-amdgcn-target.ll
@@ -18,11 +18,11 @@
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=bonaire < %s | FileCheck --check-prefixes=GFX704 %s
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx705 < %s | FileCheck --check-prefixes=GFX705 %s
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx801 < %s | FileCheck --check-prefixes=GFX801 %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx801 -mattr=-xnack < %s | FileCheck --check-prefixes=GFX801-NOXNACK %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx801 -mattr=+xnack < %s | FileCheck --check-prefixes=GFX801-XNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx801 -amdgpu-xnack=0 < %s | FileCheck --check-prefixes=GFX801-NOXNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx801 -amdgpu-xnack=1 < %s | FileCheck --check-prefixes=GFX801-XNACK %s
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=carrizo < %s | FileCheck --check-prefixes=GFX801 %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=carrizo -mattr=-xnack < %s | FileCheck --check-prefixes=GFX801-NOXNACK %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=carrizo -mattr=+xnack < %s | FileCheck --check-prefixes=GFX801-XNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=carrizo -amdgpu-xnack=0 < %s | FileCheck --check-prefixes=GFX801-NOXNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=carrizo -amdgpu-xnack=1 < %s | FileCheck --check-prefixes=GFX801-XNACK %s
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx802 < %s | FileCheck --check-prefixes=GFX802 %s
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=iceland < %s | FileCheck --check-prefixes=GFX802 %s
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=tonga < %s | FileCheck --check-prefixes=GFX802 %s
@@ -33,62 +33,62 @@
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx805 < %s | FileCheck --check-prefixes=GFX805 %s
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=tongapro < %s | FileCheck --check-prefixes=GFX805 %s
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx810 < %s | FileCheck --check-prefixes=GFX810 %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx810 -mattr=-xnack < %s | FileCheck --check-prefixes=GFX810-NOXNACK %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx810 -mattr=+xnack < %s | FileCheck --check-prefixes=GFX810-XNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx810 -amdgpu-xnack=0 < %s | FileCheck --check-prefixes=GFX810-NOXNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx810 -amdgpu-xnack=1 < %s | FileCheck --check-prefixes=GFX810-XNACK %s
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=stoney < %s | FileCheck --check-prefixes=GFX810 %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=stoney -mattr=-xnack < %s | FileCheck --check-prefixes=GFX810-NOXNACK %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=stoney -mattr=+xnack < %s | FileCheck --check-prefixes=GFX810-XNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=stoney -amdgpu-xnack=0 < %s | FileCheck --check-prefixes=GFX810-NOXNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=stoney -amdgpu-xnack=1 < %s | FileCheck --check-prefixes=GFX810-XNACK %s
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 < %s | FileCheck --check-prefixes=GFX900 %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -mattr=-xnack < %s | FileCheck --check-prefixes=GFX900-NOXNACK %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -mattr=+xnack < %s | FileCheck --check-prefixes=GFX900-XNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -amdgpu-xnack=0 < %s | FileCheck --check-prefixes=GFX900-NOXNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -amdgpu-xnack=1 < %s | FileCheck --check-prefixes=GFX900-XNACK %s
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx902 < %s | FileCheck --check-prefixes=GFX902 %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx902 -mattr=-xnack < %s | FileCheck --check-prefixes=GFX902-NOXNACK %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx902 -mattr=+xnack < %s | FileCheck --check-prefixes=GFX902-XNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx902 -amdgpu-xnack=0 < %s | FileCheck --check-prefixes=GFX902-NOXNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx902 -amdgpu-xnack=1 < %s | FileCheck --check-prefixes=GFX902-XNACK %s
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx904 < %s | FileCheck --check-prefixes=GFX904 %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx904 -mattr=-xnack < %s | FileCheck --check-prefixes=GFX904-NOXNACK %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx904 -mattr=+xnack < %s | FileCheck --check-prefixes=GFX904-XNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx904 -amdgpu-xnack=0 < %s | FileCheck --check-prefixes=GFX904-NOXNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx904 -amdgpu-xnack=1 < %s | FileCheck --check-prefixes=GFX904-XNACK %s
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 < %s | FileCheck --check-prefixes=GFX906 %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 -mattr=-sramecc < %s | FileCheck --check-prefixes=GFX906-NOSRAMECC %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 -mattr=+sramecc < %s | FileCheck --check-prefixes=GFX906-SRAMECC %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 -mattr=-xnack < %s | FileCheck --check-prefixes=GFX906-NOXNACK %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 -mattr=+xnack < %s | FileCheck --check-prefixes=GFX906-XNACK %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 -mattr=-sramecc,-xnack < %s | FileCheck --check-prefixes=GFX906-NOSRAMECC-NOXNACK %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 -mattr=+sramecc,-xnack < %s | FileCheck --check-prefixes=GFX906-SRAMECC-NOXNACK %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 -mattr=-sramecc,+xnack < %s | FileCheck --check-prefixes=GFX906-NOSRAMECC-XNACK %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 -mattr=+sramecc,+xnack < %s | FileCheck --check-prefixes=GFX906-SRAMECC-XNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 -amdgpu-sramecc=0 < %s | FileCheck --check-prefixes=GFX906-NOSRAMECC %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 -amdgpu-sramecc=1 < %s | FileCheck --check-prefixes=GFX906-SRAMECC %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 -amdgpu-xnack=0 < %s | FileCheck --check-prefixes=GFX906-NOXNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 -amdgpu-xnack=1 < %s | FileCheck --check-prefixes=GFX906-XNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 -amdgpu-sramecc=0 -amdgpu-xnack=0 < %s | FileCheck --check-prefixes=GFX906-NOSRAMECC-NOXNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 -amdgpu-sramecc=1 -amdgpu-xnack=0 < %s | FileCheck --check-prefixes=GFX906-SRAMECC-NOXNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 -amdgpu-sramecc=0 -amdgpu-xnack=1 < %s | FileCheck --check-prefixes=GFX906-NOSRAMECC-XNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 -amdgpu-sramecc=1 -amdgpu-xnack=1 < %s | FileCheck --check-prefixes=GFX906-SRAMECC-XNACK %s
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx908 < %s | FileCheck --check-prefixes=GFX908 %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx908 -mattr=-sramecc < %s | FileCheck --check-prefixes=GFX908-NOSRAMECC %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx908 -mattr=+sramecc < %s | FileCheck --check-prefixes=GFX908-SRAMECC %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx908 -mattr=-xnack < %s | FileCheck --check-prefixes=GFX908-NOXNACK %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx908 -mattr=+xnack < %s | FileCheck --check-prefixes=GFX908-XNACK %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx908 -mattr=-sramecc,-xnack < %s | FileCheck --check-prefixes=GFX908-NOSRAMECC-NOXNACK %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx908 -mattr=+sramecc,-xnack < %s | FileCheck --check-prefixes=GFX908-SRAMECC-NOXNACK %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx908 -mattr=-sramecc,+xnack < %s | FileCheck --check-prefixes=GFX908-NOSRAMECC-XNACK %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx908 -mattr=+sramecc,+xnack < %s | FileCheck --check-prefixes=GFX908-SRAMECC-XNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx908 -amdgpu-sramecc=0 < %s | FileCheck --check-prefixes=GFX908-NOSRAMECC %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx908 -amdgpu-sramecc=1 < %s | FileCheck --check-prefixes=GFX908-SRAMECC %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx908 -amdgpu-xnack=0 < %s | FileCheck --check-prefixes=GFX908-NOXNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx908 -amdgpu-xnack=1 < %s | FileCheck --check-prefixes=GFX908-XNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx908 -amdgpu-sramecc=0 -amdgpu-xnack=0 < %s | FileCheck --check-prefixes=GFX908-NOSRAMECC-NOXNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx908 -amdgpu-sramecc=1 -amdgpu-xnack=0 < %s | FileCheck --check-prefixes=GFX908-SRAMECC-NOXNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx908 -amdgpu-sramecc=0 -amdgpu-xnack=1 < %s | FileCheck --check-prefixes=GFX908-NOSRAMECC-XNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx908 -amdgpu-sramecc=1 -amdgpu-xnack=1 < %s | FileCheck --check-prefixes=GFX908-SRAMECC-XNACK %s
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx909 < %s | FileCheck --check-prefixes=GFX909 %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx909 -mattr=-xnack < %s | FileCheck --check-prefixes=GFX909-NOXNACK %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx909 -mattr=+xnack < %s | FileCheck --check-prefixes=GFX909-XNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx909 -amdgpu-xnack=0 < %s | FileCheck --check-prefixes=GFX909-NOXNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx909 -amdgpu-xnack=1 < %s | FileCheck --check-prefixes=GFX909-XNACK %s
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90c < %s | FileCheck --check-prefixes=GFX90C %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90c -mattr=-xnack < %s | FileCheck --check-prefixes=GFX90C-NOXNACK %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90c -mattr=+xnack < %s | FileCheck --check-prefixes=GFX90C-XNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90c -amdgpu-xnack=0 < %s | FileCheck --check-prefixes=GFX90C-NOXNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90c -amdgpu-xnack=1 < %s | FileCheck --check-prefixes=GFX90C-XNACK %s
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx942 < %s | FileCheck --check-prefixes=GFX942 %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx942 -mattr=-xnack < %s | FileCheck --check-prefixes=GFX942-NOXNACK %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx942 -mattr=+xnack < %s | FileCheck --check-prefixes=GFX942-XNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx942 -amdgpu-xnack=0 < %s | FileCheck --check-prefixes=GFX942-NOXNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx942 -amdgpu-xnack=1 < %s | FileCheck --check-prefixes=GFX942-XNACK %s
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 < %s | FileCheck --check-prefixes=GFX950 %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 -mattr=-xnack < %s | FileCheck --check-prefixes=GFX950-NOXNACK %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 -mattr=+xnack < %s | FileCheck --check-prefixes=GFX950-XNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 -amdgpu-xnack=0 < %s | FileCheck --check-prefixes=GFX950-NOXNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 -amdgpu-xnack=1 < %s | FileCheck --check-prefixes=GFX950-XNACK %s
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1010 < %s | FileCheck --check-prefixes=GFX1010 %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1010 -mattr=-xnack < %s | FileCheck --check-prefixes=GFX1010-NOXNACK %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1010 -mattr=+xnack < %s | FileCheck --check-prefixes=GFX1010-XNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1010 -amdgpu-xnack=0 < %s | FileCheck --check-prefixes=GFX1010-NOXNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1010 -amdgpu-xnack=1 < %s | FileCheck --check-prefixes=GFX1010-XNACK %s
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1011 < %s | FileCheck --check-prefixes=GFX1011 %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1011 -mattr=-xnack < %s | FileCheck --check-prefixes=GFX1011-NOXNACK %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1011 -mattr=+xnack < %s | FileCheck --check-prefixes=GFX1011-XNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1011 -amdgpu-xnack=0 < %s | FileCheck --check-prefixes=GFX1011-NOXNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1011 -amdgpu-xnack=1 < %s | FileCheck --check-prefixes=GFX1011-XNACK %s
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1012 < %s | FileCheck --check-prefixes=GFX1012 %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1012 -mattr=-xnack < %s | FileCheck --check-prefixes=GFX1012-NOXNACK %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1012 -mattr=+xnack < %s | FileCheck --check-prefixes=GFX1012-XNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1012 -amdgpu-xnack=0 < %s | FileCheck --check-prefixes=GFX1012-NOXNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1012 -amdgpu-xnack=1 < %s | FileCheck --check-prefixes=GFX1012-XNACK %s
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1013 < %s | FileCheck --check-prefixes=GFX1013 %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1013 -mattr=-xnack < %s | FileCheck --check-prefixes=GFX1013-NOXNACK %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1013 -mattr=+xnack < %s | FileCheck --check-prefixes=GFX1013-XNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1013 -amdgpu-xnack=0 < %s | FileCheck --check-prefixes=GFX1013-NOXNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1013 -amdgpu-xnack=1 < %s | FileCheck --check-prefixes=GFX1013-XNACK %s
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1030 < %s | FileCheck --check-prefixes=GFX1030 %s
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1031 < %s | FileCheck --check-prefixes=GFX1031 %s
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1032 < %s | FileCheck --check-prefixes=GFX1032 %s
@@ -113,12 +113,12 @@
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1251 < %s | FileCheck --check-prefixes=GFX1251 %s
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1310 < %s | FileCheck --check-prefixes=GFX1310 %s
-; RUN: llc --amdhsa-code-object-version=6 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx9-generic -mattr=-xnack < %s | FileCheck --check-prefixes=GFX9_GENERIC_NOXNACK %s
-; RUN: llc --amdhsa-code-object-version=6 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx9-generic -mattr=+xnack < %s | FileCheck --check-prefixes=GFX9_GENERIC_XNACK %s
-; RUN: llc --amdhsa-code-object-version=6 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx9-4-generic -mattr=-xnack < %s | FileCheck --check-prefixes=GFX9_4_GENERIC_NOXNACK %s
-; RUN: llc --amdhsa-code-object-version=6 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx9-4-generic -mattr=+xnack < %s | FileCheck --check-prefixes=GFX9_4_GENERIC_XNACK %s
-; RUN: llc --amdhsa-code-object-version=6 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx10-1-generic -mattr=-xnack < %s | FileCheck --check-prefixes=GFX10_1_GENERIC_NOXNACK %s
-; RUN: llc --amdhsa-code-object-version=6 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx10-1-generic -mattr=+xnack < %s | FileCheck --check-prefixes=GFX10_1_GENERIC_XNACK %s
+; RUN: llc --amdhsa-code-object-version=6 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx9-generic -amdgpu-xnack=0 < %s | FileCheck --check-prefixes=GFX9_GENERIC_NOXNACK %s
+; RUN: llc --amdhsa-code-object-version=6 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx9-generic -amdgpu-xnack=1 < %s | FileCheck --check-prefixes=GFX9_GENERIC_XNACK %s
+; RUN: llc --amdhsa-code-object-version=6 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx9-4-generic -amdgpu-xnack=0 < %s | FileCheck --check-prefixes=GFX9_4_GENERIC_NOXNACK %s
+; RUN: llc --amdhsa-code-object-version=6 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx9-4-generic -amdgpu-xnack=1 < %s | FileCheck --check-prefixes=GFX9_4_GENERIC_XNACK %s
+; RUN: llc --amdhsa-code-object-version=6 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx10-1-generic -amdgpu-xnack=0 < %s | FileCheck --check-prefixes=GFX10_1_GENERIC_NOXNACK %s
+; RUN: llc --amdhsa-code-object-version=6 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx10-1-generic -amdgpu-xnack=1 < %s | FileCheck --check-prefixes=GFX10_1_GENERIC_XNACK %s
; RUN: llc --amdhsa-code-object-version=6 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx10-3-generic < %s | FileCheck --check-prefixes=GFX10_3_GENERIC %s
; RUN: llc --amdhsa-code-object-version=6 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx11-generic < %s | FileCheck --check-prefixes=GFX11_GENERIC %s
; RUN: llc --amdhsa-code-object-version=6 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx12-generic < %s | FileCheck --check-prefixes=GFX12_GENERIC %s
diff --git a/llvm/test/CodeGen/AMDGPU/elf-header-flags-sramecc.ll b/llvm/test/CodeGen/AMDGPU/elf-header-flags-sramecc.ll
index 09fb3e092a307..834971c07422d 100644
--- a/llvm/test/CodeGen/AMDGPU/elf-header-flags-sramecc.ll
+++ b/llvm/test/CodeGen/AMDGPU/elf-header-flags-sramecc.ll
@@ -1,22 +1,22 @@
; RUN: llc -filetype=obj -mtriple=amdgcn -mcpu=gfx906 < %s | llvm-readobj --file-headers - | FileCheck --check-prefixes=SRAM-ECC-GFX906 %s
-; RUN: llc -filetype=obj -mtriple=amdgcn -mcpu=gfx906 -mattr=-sramecc < %s | llvm-readobj --file-headers - | FileCheck --check-prefixes=NO-SRAM-ECC-GFX906 %s
-; RUN: llc -filetype=obj -mtriple=amdgcn -mcpu=gfx906 -mattr=+sramecc < %s | llvm-readobj --file-headers - | FileCheck --check-prefixes=SRAM-ECC-GFX906 %s
-; RUN: llc -filetype=obj -mtriple=amdgcn -mcpu=gfx906 -mattr=+sramecc,+xnack < %s | llvm-readobj --file-headers - | FileCheck --check-prefixes=SRAM-ECC-XNACK-GFX906 %s
+; RUN: llc -filetype=obj -mtriple=amdgcn -mcpu=gfx906 -amdgpu-sramecc=0 < %s | llvm-readobj --file-headers - | FileCheck --check-prefixes=NO-SRAM-ECC-GFX906 %s
+; RUN: llc -filetype=obj -mtriple=amdgcn -mcpu=gfx906 -amdgpu-sramecc=1 < %s | llvm-readobj --file-headers - | FileCheck --check-prefixes=SRAM-ECC-GFX906 %s
+; RUN: llc -filetype=obj -mtriple=amdgcn -mcpu=gfx906 -amdgpu-sramecc=1 -amdgpu-xnack=1 < %s | llvm-readobj --file-headers - | FileCheck --check-prefixes=SRAM-ECC-XNACK-GFX906 %s
; RUN: llc -filetype=obj -mtriple=amdgcn -mcpu=gfx908 < %s | llvm-readobj --file-header - | FileCheck --check-prefix=SRAM-ECC-GFX908 %s
-; RUN: llc -filetype=obj -mtriple=amdgcn -mcpu=gfx908 -mattr=+sramecc < %s | llvm-readobj --file-header - | FileCheck --check-prefix=SRAM-ECC-GFX908 %s
+; RUN: llc -filetype=obj -mtriple=amdgcn -mcpu=gfx908 -amdgpu-sramecc=1 < %s | llvm-readobj --file-header - | FileCheck --check-prefix=SRAM-ECC-GFX908 %s
; RUN: llc -filetype=obj -mtriple=amdgcn -mcpu=gfx90a < %s | llvm-readobj --file-header - | FileCheck --check-prefix=SRAM-ECC-GFX90A %s
-; RUN: llc -filetype=obj -mtriple=amdgcn -mcpu=gfx90a -mattr=+sramecc < %s | llvm-readobj --file-header - | FileCheck --check-prefix=SRAM-ECC-GFX90A %s
+; RUN: llc -filetype=obj -mtriple=amdgcn -mcpu=gfx90a -amdgpu-sramecc=1 < %s | llvm-readobj --file-header - | FileCheck --check-prefix=SRAM-ECC-GFX90A %s
; RUN: llc -filetype=obj -mtriple=amdgcn -mcpu=gfx942 < %s | llvm-readobj --file-header - | FileCheck --check-prefix=SRAM-ECC-GFX942 %s
-; RUN: llc -filetype=obj -mtriple=amdgcn -mcpu=gfx942 -mattr=+sramecc < %s | llvm-readobj --file-header - | FileCheck --check-prefix=SRAM-ECC-GFX942 %s
+; RUN: llc -filetype=obj -mtriple=amdgcn -mcpu=gfx942 -amdgpu-sramecc=1 < %s | llvm-readobj --file-header - | FileCheck --check-prefix=SRAM-ECC-GFX942 %s
; RUN: llc -filetype=obj -mtriple=amdgcn -mcpu=gfx950 < %s | llvm-readobj --file-header - | FileCheck --check-prefix=SRAM-ECC-GFX950 %s
-; RUN: llc -filetype=obj -mtriple=amdgcn -mcpu=gfx950 -mattr=+sramecc < %s | llvm-readobj --file-header - | FileCheck --check-prefix=SRAM-ECC-GFX950 %s
+; RUN: llc -filetype=obj -mtriple=amdgcn -mcpu=gfx950 -amdgpu-sramecc=1 < %s | llvm-readobj --file-header - | FileCheck --check-prefix=SRAM-ECC-GFX950 %s
; RUN: llc -filetype=obj -mtriple=amdgcn -mcpu=gfx1250 < %s | llvm-readobj --file-header - | FileCheck --check-prefix=SRAM-ECC-GFX1250 %s
-; RUN: llc -filetype=obj -mtriple=amdgcn -mcpu=gfx1250 -mattr=+sramecc < %s | llvm-readobj --file-header - | FileCheck --check-prefix=SRAM-ECC-GFX1250 %s
+; RUN: llc -filetype=obj -mtriple=amdgcn -mcpu=gfx1250 -amdgpu-sramecc=1 < %s | llvm-readobj --file-header - | FileCheck --check-prefix=SRAM-ECC-GFX1250 %s
; NO-SRAM-ECC-GFX906: Flags [
; NO-SRAM-ECC-GFX906-NEXT: EF_AMDGPU_FEATURE_XNACK_V3 (0x100)
diff --git a/llvm/test/CodeGen/AMDGPU/elf-header-flags-xnack.ll b/llvm/test/CodeGen/AMDGPU/elf-header-flags-xnack.ll
index 98aae6d603eda..a9716e57a331b 100644
--- a/llvm/test/CodeGen/AMDGPU/elf-header-flags-xnack.ll
+++ b/llvm/test/CodeGen/AMDGPU/elf-header-flags-xnack.ll
@@ -1,7 +1,7 @@
; RUN: llc -filetype=obj -mtriple=amdgcn -mcpu=gfx801 < %s | llvm-readobj --file-headers - | FileCheck --check-prefixes=XNACK-GFX801 %s
-; RUN: llc -filetype=obj -mtriple=amdgcn -mcpu=gfx801 -mattr=+xnack < %s | llvm-readobj --file-headers - | FileCheck --check-prefixes=XNACK-GFX801 %s
+; RUN: llc -filetype=obj -mtriple=amdgcn -mcpu=gfx801 --amdgpu-xnack=true < %s | llvm-readobj --file-headers - | FileCheck --check-prefixes=XNACK-GFX801 %s
; RUN: llc -filetype=obj -mtriple=amdgcn -mcpu=gfx802 < %s | llvm-readobj --file-headers - | FileCheck --check-prefixes=NO-XNACK-GFX802 %s
-; RUN: llc -filetype=obj -mtriple=amdgcn -mcpu=gfx802 -mattr=-xnack < %s | llvm-readobj --file-headers - | FileCheck --check-prefixes=NO-XNACK-GFX802 %s
+; RUN: llc -filetype=obj -mtriple=amdgcn -mcpu=gfx802 --amdgpu-xnack=false < %s | llvm-readobj --file-headers - | FileCheck --check-prefixes=NO-XNACK-GFX802 %s
; XNACK-GFX801: Flags [
; XNACK-GFX801-NEXT: EF_AMDGPU_FEATURE_XNACK_V3 (0x100)
diff --git a/llvm/test/CodeGen/AMDGPU/flat-saddr-load.ll b/llvm/test/CodeGen/AMDGPU/flat-saddr-load.ll
index 6ac69cc7e2f51..9dc26de7a431b 100644
--- a/llvm/test/CodeGen/AMDGPU/flat-saddr-load.ll
+++ b/llvm/test/CodeGen/AMDGPU/flat-saddr-load.ll
@@ -4,8 +4,8 @@
; RUN: llc -global-isel=0 -mtriple=amdgcn-mesa-mesa3d -mcpu=gfx1250 -mattr=+real-true16 < %s | FileCheck -check-prefixes=GFX1250,GFX1250-SDAG,GFX1250-SDAG-TRUE16 %s
; RUN: llc -global-isel=1 -new-reg-bank-select -mtriple=amdgcn-mesa-mesa3d -mcpu=gfx1250 -mattr=+real-true16 < %s | FileCheck -check-prefixes=GFX1250,GFX1250-GISEL,GFX1250-GISEL-TRUE16 %s
-; RUN: llc -mtriple=amdgcn-mesa-mesa3d -mcpu=gfx1250 -mattr=+real-true16,-sramecc < %s | FileCheck -check-prefixes=GFX1250,GFX1250-NOECC,GFX1250-NOECC-SDAG-TRUE16 %s
-; RUN: llc -mtriple=amdgcn-mesa-mesa3d -mcpu=gfx1250 -mattr=-real-true16,-sramecc < %s | FileCheck -check-prefixes=GFX1250,GFX1250-NOECC,GFX1250-NOECC-SDAG-FAKE16 %s
+; RUN: llc -mtriple=amdgcn-mesa-mesa3d -mcpu=gfx1250 -mattr=+real-true16 --amdgpu-sramecc=false < %s | FileCheck -check-prefixes=GFX1250,GFX1250-NOECC,GFX1250-NOECC-SDAG-TRUE16 %s
+; RUN: llc -mtriple=amdgcn-mesa-mesa3d -mcpu=gfx1250 -mattr=-real-true16 --amdgpu-sramecc=false < %s | FileCheck -check-prefixes=GFX1250,GFX1250-NOECC,GFX1250-NOECC-SDAG-FAKE16 %s
; Test using saddr addressing mode of flat_*load_* instructions.
diff --git a/llvm/test/CodeGen/AMDGPU/flat-scratch-reg.ll b/llvm/test/CodeGen/AMDGPU/flat-scratch-reg.ll
index a59382ba20dc5..30389375292b3 100644
--- a/llvm/test/CodeGen/AMDGPU/flat-scratch-reg.ll
+++ b/llvm/test/CodeGen/AMDGPU/flat-scratch-reg.ll
@@ -1,23 +1,23 @@
; RUN: llc < %s -mtriple=amdgcn -mcpu=kaveri | FileCheck -check-prefix=CI -check-prefix=GCN %s
-; RUN: llc < %s -mtriple=amdgcn -mcpu=fiji -mattr=-xnack | FileCheck -check-prefix=VI-NOXNACK -check-prefix=GCN %s
+; RUN: llc < %s -mtriple=amdgcn -mcpu=fiji --amdgpu-xnack=false | FileCheck -check-prefix=VI-NOXNACK -check-prefix=GCN %s
-; RUN: llc < %s -mtriple=amdgcn -mcpu=carrizo -mattr=-xnack | FileCheck -check-prefixes=VI-NOXNACK,GCN %s
-; RUN: llc < %s -mtriple=amdgcn -mcpu=stoney -mattr=-xnack | FileCheck -check-prefixes=VI-NOXNACK,GCN %s
+; RUN: llc < %s -mtriple=amdgcn -mcpu=carrizo --amdgpu-xnack=false | FileCheck -check-prefixes=VI-NOXNACK,GCN %s
+; RUN: llc < %s -mtriple=amdgcn -mcpu=stoney --amdgpu-xnack=false | FileCheck -check-prefixes=VI-NOXNACK,GCN %s
-; RUN: llc < %s -mtriple=amdgcn -mcpu=carrizo -mattr=+xnack | FileCheck -check-prefix=VI-XNACK -check-prefix=GCN %s
-; RUN: llc < %s -mtriple=amdgcn -mcpu=stoney -mattr=+xnack | FileCheck -check-prefix=VI-XNACK -check-prefix=GCN %s
+; RUN: llc < %s -mtriple=amdgcn -mcpu=carrizo --amdgpu-xnack=true | FileCheck -check-prefix=VI-XNACK -check-prefix=GCN %s
+; RUN: llc < %s -mtriple=amdgcn -mcpu=stoney --amdgpu-xnack=true | FileCheck -check-prefix=VI-XNACK -check-prefix=GCN %s
; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=kaveri | FileCheck -check-prefixes=GCN %s
-; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=carrizo -mattr=-xnack | FileCheck -check-prefixes=HSA-VI-NOXNACK,GCN %s
-; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=carrizo -mattr=+xnack | FileCheck -check-prefixes=HSA-VI-XNACK,GCN %s
+; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=carrizo --amdgpu-xnack=false | FileCheck -check-prefixes=HSA-VI-NOXNACK,GCN %s
+; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=carrizo --amdgpu-xnack=true | FileCheck -check-prefixes=HSA-VI-XNACK,GCN %s
; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=gfx900 -mattr=+architected-flat-scratch | FileCheck -check-prefixes=GCN %s
-; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=gfx900 -mattr=+architected-flat-scratch,-xnack | FileCheck -check-prefixes=GFX9-ARCH-FLAT-NOXNACK,GCN %s
-; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=gfx900 -mattr=+architected-flat-scratch,+xnack | FileCheck -check-prefixes=GFX9-ARCH-FLAT-XNACK,GCN %s
+; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=gfx900 -mattr=+architected-flat-scratch --amdgpu-xnack=false | FileCheck -check-prefixes=GFX9-ARCH-FLAT-NOXNACK,GCN %s
+; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=gfx900 -mattr=+architected-flat-scratch --amdgpu-xnack=true | FileCheck -check-prefixes=GFX9-ARCH-FLAT-XNACK,GCN %s
; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=gfx1010 -mattr=+architected-flat-scratch | FileCheck -check-prefixes=GCN %s
-; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=gfx1010 -mattr=+architected-flat-scratch,-xnack | FileCheck -check-prefixes=GFX10-ARCH-FLAT-NOXNACK,GCN %s
-; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=gfx1010 -mattr=+architected-flat-scratch,+xnack | FileCheck -check-prefixes=GFX10-ARCH-FLAT-XNACK,GCN %s
+; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=gfx1010 -mattr=+architected-flat-scratch --amdgpu-xnack=false | FileCheck -check-prefixes=GFX10-ARCH-FLAT-NOXNACK,GCN %s
+; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=gfx1010 -mattr=+architected-flat-scratch --amdgpu-xnack=true | FileCheck -check-prefixes=GFX10-ARCH-FLAT-XNACK,GCN %s
; GCN-LABEL: {{^}}no_vcc_no_flat:
diff --git a/llvm/test/CodeGen/AMDGPU/gfx902-without-xnack.ll b/llvm/test/CodeGen/AMDGPU/gfx902-without-xnack.ll
index a83cde14892b5..fcf828241303a 100644
--- a/llvm/test/CodeGen/AMDGPU/gfx902-without-xnack.ll
+++ b/llvm/test/CodeGen/AMDGPU/gfx902-without-xnack.ll
@@ -1,4 +1,4 @@
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx902 -mattr=-xnack < %s | FileCheck %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx902 < %s | FileCheck %s
; CHECK: .amdgcn_target "amdgcn-amd-amdhsa--gfx902:xnack-"
define amdgpu_kernel void @test_kernel(ptr addrspace(1) %out0, ptr addrspace(1) %out1) nounwind {
@@ -6,5 +6,6 @@ define amdgpu_kernel void @test_kernel(ptr addrspace(1) %out0, ptr addrspace(1)
ret void
}
-!llvm.module.flags = !{!0}
+!llvm.module.flags = !{!0, !1}
!0 = !{i32 1, !"amdhsa_code_object_version", i32 400}
+!1 = !{i32 1, !"amdgpu.xnack", i32 0}
diff --git a/llvm/test/CodeGen/AMDGPU/greedy-reverse-local-assignment.ll b/llvm/test/CodeGen/AMDGPU/greedy-reverse-local-assignment.ll
index 6e82a294243d2..5eda20acd4604 100644
--- a/llvm/test/CodeGen/AMDGPU/greedy-reverse-local-assignment.ll
+++ b/llvm/test/CodeGen/AMDGPU/greedy-reverse-local-assignment.ll
@@ -2,8 +2,8 @@
; RUN: llc -mtriple=amdgcn-amd-amdhsa -greedy-reverse-local-assignment=0 -mcpu=gfx900 < %s | FileCheck -check-prefixes=FORWARDXNACK %s
; RUN: llc -mtriple=amdgcn-amd-amdhsa -greedy-reverse-local-assignment=1 -mcpu=gfx900 < %s | FileCheck -check-prefixes=REVERSEXNACK %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -greedy-reverse-local-assignment=0 -mcpu=gfx900 -mattr=-xnack < %s | FileCheck -check-prefix=NOXNACK %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -greedy-reverse-local-assignment=1 -mcpu=gfx900 -mattr=-xnack < %s | FileCheck -check-prefix=NOXNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -greedy-reverse-local-assignment=0 -mcpu=gfx900 --amdgpu-xnack=false < %s | FileCheck -check-prefix=NOXNACK %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -greedy-reverse-local-assignment=1 -mcpu=gfx900 --amdgpu-xnack=false < %s | FileCheck -check-prefix=NOXNACK %s
; Test the change in the behavior of the allocator with
; -greedy-reverse-local-reassignment enabled. This case shows a
diff --git a/llvm/test/CodeGen/AMDGPU/hazard-hidden-bundle.mir b/llvm/test/CodeGen/AMDGPU/hazard-hidden-bundle.mir
index e2635fc6ecbbb..757aba9587666 100644
--- a/llvm/test/CodeGen/AMDGPU/hazard-hidden-bundle.mir
+++ b/llvm/test/CodeGen/AMDGPU/hazard-hidden-bundle.mir
@@ -1,6 +1,6 @@
-# RUN: llc -mtriple=amdgcn -mcpu=gfx902 -mattr=+xnack -verify-machineinstrs -run-pass post-RA-hazard-rec %s -o - | FileCheck -check-prefixes=GCN,XNACK,GFX9 %s
-# RUN: llc -mtriple=amdgcn -mcpu=gfx900 -mattr=-xnack -verify-machineinstrs -run-pass post-RA-hazard-rec %s -o - | FileCheck -check-prefixes=GCN,NOXNACK,GFX9 %s
-# RUN: llc -mtriple=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize64,-xnack -verify-machineinstrs -run-pass post-RA-hazard-rec %s -o - | FileCheck -check-prefixes=GCN,NOXNACK %s
+# RUN: llc -mtriple=amdgcn -mcpu=gfx902 --amdgpu-xnack=true -verify-machineinstrs -run-pass post-RA-hazard-rec %s -o - | FileCheck -check-prefixes=GCN,XNACK,GFX9 %s
+# RUN: llc -mtriple=amdgcn -mcpu=gfx900 --amdgpu-xnack=false -verify-machineinstrs -run-pass post-RA-hazard-rec %s -o - | FileCheck -check-prefixes=GCN,NOXNACK,GFX9 %s
+# RUN: llc -mtriple=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize64 --amdgpu-xnack=false -verify-machineinstrs -run-pass post-RA-hazard-rec %s -o - | FileCheck -check-prefixes=GCN,NOXNACK %s
# GCN-LABEL: name: break_smem_clause_simple_load_smrd8_ptr_hidden_bundle
# GCN: bb.0:
diff --git a/llvm/test/CodeGen/AMDGPU/hazard-in-bundle.mir b/llvm/test/CodeGen/AMDGPU/hazard-in-bundle.mir
index 401f6e303e796..349aba9b8140f 100644
--- a/llvm/test/CodeGen/AMDGPU/hazard-in-bundle.mir
+++ b/llvm/test/CodeGen/AMDGPU/hazard-in-bundle.mir
@@ -1,6 +1,6 @@
-# RUN: llc -mtriple=amdgcn -mcpu=gfx902 -mattr=+xnack -verify-machineinstrs -run-pass post-RA-hazard-rec %s -o - | FileCheck -check-prefixes=GCN,XNACK,GFX9 %s
-# RUN: llc -mtriple=amdgcn -mcpu=gfx900 -mattr=-xnack -verify-machineinstrs -run-pass post-RA-hazard-rec %s -o - | FileCheck -check-prefixes=GCN,NOXNACK,GFX9 %s
-# RUN: llc -mtriple=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize64,-xnack -verify-machineinstrs -run-pass post-RA-hazard-rec %s -o - | FileCheck -check-prefixes=GCN,NOXNACK,GFX10 %s
+# RUN: llc -mtriple=amdgcn -mcpu=gfx902 --amdgpu-xnack=true -verify-machineinstrs -run-pass post-RA-hazard-rec %s -o - | FileCheck -check-prefixes=GCN,XNACK,GFX9 %s
+# RUN: llc -mtriple=amdgcn -mcpu=gfx900 --amdgpu-xnack=false -verify-machineinstrs -run-pass post-RA-hazard-rec %s -o - | FileCheck -check-prefixes=GCN,NOXNACK,GFX9 %s
+# RUN: llc -mtriple=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize64 --amdgpu-xnack=false -verify-machineinstrs -run-pass post-RA-hazard-rec %s -o - | FileCheck -check-prefixes=GCN,NOXNACK,GFX10 %s
# GCN-LABEL: name: break_smem_clause_max_look_ahead_in_bundle
# GCN: S_LOAD_DWORDX2_IMM
diff --git a/llvm/test/CodeGen/AMDGPU/hsa-metadata-kernel-code-props.ll b/llvm/test/CodeGen/AMDGPU/hsa-metadata-kernel-code-props.ll
index cd89a36fe538b..fdec81ed59f03 100644
--- a/llvm/test/CodeGen/AMDGPU/hsa-metadata-kernel-code-props.ll
+++ b/llvm/test/CodeGen/AMDGPU/hsa-metadata-kernel-code-props.ll
@@ -1,7 +1,7 @@
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx700 -enable-misched=0 -filetype=obj -o - < %s | llvm-readelf --notes - | FileCheck --check-prefixes=CHECK,GFX700,WAVE64 %s
-; RUN: llc -mattr=-xnack -mtriple=amdgcn-amd-amdhsa -mcpu=gfx803 -enable-misched=0 -filetype=obj -o - < %s | llvm-readelf --notes - | FileCheck --check-prefixes=CHECK,GFX803,WAVE64 %s
-; RUN: llc -mattr=-xnack -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -enable-misched=0 -filetype=obj -o - < %s | llvm-readelf --notes - | FileCheck --check-prefixes=CHECK,GFX900,WAVE64 %s
-; RUN: llc -mattr=-xnack -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1010 -enable-misched=0 -filetype=obj -o - < %s | llvm-readelf --notes - | FileCheck --check-prefixes=CHECK,GFX1010,WAVE32 %s
+; RUN: llc --amdgpu-xnack=false -mtriple=amdgcn-amd-amdhsa -mcpu=gfx803 -enable-misched=0 -filetype=obj -o - < %s | llvm-readelf --notes - | FileCheck --check-prefixes=CHECK,GFX803,WAVE64 %s
+; RUN: llc --amdgpu-xnack=false -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -enable-misched=0 -filetype=obj -o - < %s | llvm-readelf --notes - | FileCheck --check-prefixes=CHECK,GFX900,WAVE64 %s
+; RUN: llc --amdgpu-xnack=false -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1010 -enable-misched=0 -filetype=obj -o - < %s | llvm-readelf --notes - | FileCheck --check-prefixes=CHECK,GFX1010,WAVE32 %s
@var = addrspace(1) global float 0.0
diff --git a/llvm/test/CodeGen/AMDGPU/hsa-metadata-resource-usage-function-ordering.ll b/llvm/test/CodeGen/AMDGPU/hsa-metadata-resource-usage-function-ordering.ll
index 03242b69beb8c..f5588c0c4ab7b 100644
--- a/llvm/test/CodeGen/AMDGPU/hsa-metadata-resource-usage-function-ordering.ll
+++ b/llvm/test/CodeGen/AMDGPU/hsa-metadata-resource-usage-function-ordering.ll
@@ -2,9 +2,9 @@
; test assertions are unlikely to succeed by accident.
; RUN: llc -amdgpu-assume-external-call-stack-size=5310 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx700 -enable-misched=0 -filetype=asm -o - < %s | FileCheck --check-prefixes CHECK,GFX7 %s
-; RUN: llc -amdgpu-assume-external-call-stack-size=5310 -mattr=-xnack -mtriple=amdgcn-amd-amdhsa -mcpu=gfx803 -enable-misched=0 -filetype=asm -o - < %s | FileCheck --check-prefixes CHECK,GFX8 %s
-; RUN: llc -amdgpu-assume-external-call-stack-size=5310 -mattr=-xnack -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -enable-misched=0 -filetype=asm -o - < %s | FileCheck --check-prefixes CHECK,GFX9 %s
-; RUN: llc -amdgpu-assume-external-call-stack-size=5310 -mattr=-xnack -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1010 -enable-misched=0 -filetype=asm -o - < %s | FileCheck --check-prefixes CHECK,GFX10 %s
+; RUN: llc -amdgpu-assume-external-call-stack-size=5310 --amdgpu-xnack=false -mtriple=amdgcn-amd-amdhsa -mcpu=gfx803 -enable-misched=0 -filetype=asm -o - < %s | FileCheck --check-prefixes CHECK,GFX8 %s
+; RUN: llc -amdgpu-assume-external-call-stack-size=5310 --amdgpu-xnack=false -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -enable-misched=0 -filetype=asm -o - < %s | FileCheck --check-prefixes CHECK,GFX9 %s
+; RUN: llc -amdgpu-assume-external-call-stack-size=5310 --amdgpu-xnack=false -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1010 -enable-misched=0 -filetype=asm -o - < %s | FileCheck --check-prefixes CHECK,GFX10 %s
; CHECK-LABEL: amdhsa.kernels
diff --git a/llvm/test/CodeGen/AMDGPU/hsa-note-no-func.ll b/llvm/test/CodeGen/AMDGPU/hsa-note-no-func.ll
index ea578fc64c699..1c628a9a14d80 100644
--- a/llvm/test/CodeGen/AMDGPU/hsa-note-no-func.ll
+++ b/llvm/test/CodeGen/AMDGPU/hsa-note-no-func.ll
@@ -25,13 +25,13 @@
; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=tongapro | FileCheck --check-prefix=HSA-VI805 %s
; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=gfx810 | FileCheck --check-prefix=HSA-VI810 %s
; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=stoney | FileCheck --check-prefix=HSA-VI810 %s
-; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=gfx900 -mattr=-xnack | FileCheck --check-prefix=HSA-GFX900 %s
+; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=gfx900 --amdgpu-xnack=false | FileCheck --check-prefix=HSA-GFX900 %s
; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=gfx900 | FileCheck --check-prefix=HSA-GFX901 %s
-; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=gfx902 -mattr=-xnack | FileCheck --check-prefix=HSA-GFX902 %s
+; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=gfx902 --amdgpu-xnack=false | FileCheck --check-prefix=HSA-GFX902 %s
; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=gfx902 | FileCheck --check-prefix=HSA-GFX903 %s
-; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=gfx904 -mattr=-xnack | FileCheck --check-prefix=HSA-GFX904 %s
+; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=gfx904 --amdgpu-xnack=false | FileCheck --check-prefix=HSA-GFX904 %s
; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=gfx904 | FileCheck --check-prefix=HSA-GFX905 %s
-; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=gfx906 -mattr=-xnack | FileCheck --check-prefix=HSA-GFX906 %s
+; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=gfx906 --amdgpu-xnack=false | FileCheck --check-prefix=HSA-GFX906 %s
; RUN: llc < %s -mtriple=amdgcn--amdhsa -mcpu=gfx906 | FileCheck --check-prefix=HSA-GFX907 %s
; NONHSA-SI600: .amd_amdgpu_isa "amdgcn-unknown-unknown--gfx600"
diff --git a/llvm/test/CodeGen/AMDGPU/immv216.ll b/llvm/test/CodeGen/AMDGPU/immv216.ll
index ae5199a3f7906..0c3dfb87e3eef 100644
--- a/llvm/test/CodeGen/AMDGPU/immv216.ll
+++ b/llvm/test/CodeGen/AMDGPU/immv216.ll
@@ -1,7 +1,7 @@
; RUN: llc -amdgpu-scalarize-global-loads=false -mtriple=amdgcn--amdhsa -mcpu=gfx1100 -mattr=-flat-for-global -show-mc-encoding < %s | FileCheck -enable-var-scope -check-prefixes=GCN,GFX10 %s
-; RUN: llc -amdgpu-scalarize-global-loads=false -mtriple=amdgcn--amdhsa -mcpu=gfx1010 -mattr=-flat-for-global,-xnack -show-mc-encoding < %s | FileCheck -enable-var-scope -check-prefixes=GCN,GFX10 %s
-; RUN: llc -amdgpu-scalarize-global-loads=false -mtriple=amdgcn--amdhsa -mcpu=gfx900 -mattr=-flat-for-global,-xnack -show-mc-encoding < %s | FileCheck -enable-var-scope -check-prefixes=GCN,GFX9 %s
-; RUN: llc -amdgpu-scalarize-global-loads=false -mtriple=amdgcn--amdhsa -mcpu=fiji -mattr=-flat-for-global,-xnack -show-mc-encoding < %s | FileCheck -enable-var-scope -check-prefixes=GCN,VI %s
+; RUN: llc -amdgpu-scalarize-global-loads=false -mtriple=amdgcn--amdhsa -mcpu=gfx1010 -mattr=-flat-for-global --amdgpu-xnack=false -show-mc-encoding < %s | FileCheck -enable-var-scope -check-prefixes=GCN,GFX10 %s
+; RUN: llc -amdgpu-scalarize-global-loads=false -mtriple=amdgcn--amdhsa -mcpu=gfx900 -mattr=-flat-for-global --amdgpu-xnack=false -show-mc-encoding < %s | FileCheck -enable-var-scope -check-prefixes=GCN,GFX9 %s
+; RUN: llc -amdgpu-scalarize-global-loads=false -mtriple=amdgcn--amdhsa -mcpu=fiji -mattr=-flat-for-global --amdgpu-xnack=false -show-mc-encoding < %s | FileCheck -enable-var-scope -check-prefixes=GCN,VI %s
; RUN: llc -amdgpu-scalarize-global-loads=false -mtriple=amdgcn--amdhsa -mcpu=kaveri -mattr=-flat-for-global -show-mc-encoding < %s | FileCheck -enable-var-scope -check-prefixes=GCN %s
; FIXME: Merge into imm.ll
diff --git a/llvm/test/CodeGen/AMDGPU/limit-soft-clause-reg-pressure.mir b/llvm/test/CodeGen/AMDGPU/limit-soft-clause-reg-pressure.mir
index bd46754d10683..9b0060277e0f1 100644
--- a/llvm/test/CodeGen/AMDGPU/limit-soft-clause-reg-pressure.mir
+++ b/llvm/test/CodeGen/AMDGPU/limit-soft-clause-reg-pressure.mir
@@ -1,5 +1,5 @@
-# RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 -mattr=+xnack -run-pass=si-form-memory-clauses -verify-machineinstrs -o - %s | FileCheck %s
-# RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 -mattr=+xnack -passes="si-form-memory-clauses" -o - %s | FileCheck %s
+# RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 --amdgpu-xnack=true -run-pass=si-form-memory-clauses -verify-machineinstrs -o - %s | FileCheck %s
+# RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 --amdgpu-xnack=true -passes="si-form-memory-clauses" -o - %s | FileCheck %s
# This previously would produce a bundle that could not be satisfied
# due to using nearly the entire register budget and not considering
diff --git a/llvm/test/CodeGen/AMDGPU/materialize-frame-index-sgpr.ll b/llvm/test/CodeGen/AMDGPU/materialize-frame-index-sgpr.ll
index edf020cce0fcc..81c3ddb6a31e4 100644
--- a/llvm/test/CodeGen/AMDGPU/materialize-frame-index-sgpr.ll
+++ b/llvm/test/CodeGen/AMDGPU/materialize-frame-index-sgpr.ll
@@ -1,8 +1,8 @@
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx700 < %s | FileCheck -check-prefix=GFX7 %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx810 -mattr=+xnack < %s | FileCheck -check-prefix=GFX8 %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -mattr=+xnack < %s | FileCheck -check-prefixes=GFX900 %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx942 -mattr=+xnack < %s | FileCheck -check-prefixes=GFX942 %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx810 --amdgpu-xnack=true < %s | FileCheck -check-prefix=GFX8 %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 --amdgpu-xnack=true < %s | FileCheck -check-prefixes=GFX900 %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx942 --amdgpu-xnack=true < %s | FileCheck -check-prefixes=GFX942 %s
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1010 < %s | FileCheck -check-prefix=GFX10_1 %s
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1030 < %s | FileCheck -check-prefix=GFX10_3 %s
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1100 < %s | FileCheck -check-prefix=GFX11 %s
diff --git a/llvm/test/CodeGen/AMDGPU/mattr-xnack-sramecc-legacy.ll b/llvm/test/CodeGen/AMDGPU/mattr-xnack-sramecc-legacy.ll
new file mode 100644
index 0000000000000..d208296868c15
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/mattr-xnack-sramecc-legacy.ll
@@ -0,0 +1,23 @@
+; Test that -mattr=±xnack/±sramecc emit errors in codegen
+; xnack/sramecc should be specified via module flags instead of subtarget features.
+;
+; RUN: not llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a -mattr=+xnack,+sramecc < %s 2>&1 | FileCheck --check-prefix=BOTH-ERR %s
+; RUN: not llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a -mattr=+xnack < %s 2>&1 | FileCheck --check-prefix=XNACK-ERR %s
+; RUN: not llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a -mattr=-xnack < %s 2>&1 | FileCheck --check-prefix=XNACK-ERR %s
+; RUN: not llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a -mattr=xnack < %s 2>&1 | FileCheck --check-prefix=XNACK-ERR %s
+; RUN: not llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a -mattr=+sramecc < %s 2>&1 | FileCheck --check-prefix=SRAMECC-ERR %s
+; RUN: not llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a -mattr=-sramecc < %s 2>&1 | FileCheck --check-prefix=SRAMECC-ERR %s
+; RUN: not llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a -mattr=sramecc < %s 2>&1 | FileCheck --check-prefix=SRAMECC-ERR %s
+
+; BOTH-ERR: error: xnack/sramecc should be specified via module flags. Use module flag 'amdgpu.xnack' instead of subtarget feature
+; BOTH-ERR: error: xnack/sramecc should be specified via module flags. Use module flag 'amdgpu.sramecc' instead of subtarget feature
+
+; XNACK-ERR: error: xnack/sramecc should be specified via module flags. Use module flag 'amdgpu.xnack' instead of subtarget feature
+; XNACK-ERR-NOT: sramecc
+
+; SRAMECC-ERR: error: xnack/sramecc should be specified via module flags. Use module flag 'amdgpu.sramecc' instead of subtarget feature
+; SRAMECC-ERR-NOT: xnack
+
+define void @kernel() {
+ ret void
+}
diff --git a/llvm/test/CodeGen/AMDGPU/module-flag-sramecc.ll b/llvm/test/CodeGen/AMDGPU/module-flag-sramecc.ll
new file mode 100644
index 0000000000000..a4b4ceac2ca88
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/module-flag-sramecc.ll
@@ -0,0 +1,66 @@
+; Test that sramecc settings are controlled by the amdgpu.sramecc
+; module flag.
+
+; RUN: split-file %s %t
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 < %t/on.ll | FileCheck --check-prefix=SRAMECC-ON %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 < %t/off.ll | FileCheck --check-prefix=SRAMECC-OFF %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 < %t/absent.ll | FileCheck --check-prefix=SRAMECC-ANY %s
+
+; Test that the is ignored on targets that don't support it. gfx906 supports sramecc, gfx900 does not.
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 < %t/on.ll | FileCheck --check-prefix=GFX900 %s
+
+; Target directives for supported target
+; SRAMECC-ON: .amdgcn_target "amdgcn-amd-amdhsa--gfx906:sramecc+"
+; SRAMECC-OFF: .amdgcn_target "amdgcn-amd-amdhsa--gfx906:sramecc-"
+; SRAMECC-ANY: .amdgcn_target "amdgcn-amd-amdhsa--gfx906"
+
+; Unsupported target ignores the flag
+; GFX900: .amdgcn_target "amdgcn-amd-amdhsa--gfx900"
+
+; When sramecc is on, avoid _d16_hi
+; SRAMECC-ON-LABEL: {{^}}load_d16:
+; SRAMECC-ON: s_waitcnt
+; SRAMECC-ON: flat_load_ushort v{{[0-9]+}}, v[{{[0-9:]+}}]
+; SRAMECC-ON: s_setpc_b64
+
+; When sramecc is off, use _d16_hi instructions
+; SRAMECC-OFF-LABEL: {{^}}load_d16:
+; SRAMECC-OFF: s_waitcnt
+; SRAMECC-OFF: flat_load_short_d16_hi v0, v[{{[0-9:]+}}]
+; SRAMECC-OFF: s_setpc_b64
+
+; SRAMECC-ANY-LABEL: {{^}}load_d16:
+; SRAMECC-ANY: s_waitcnt
+; SRAMECC-ANY: flat_load_ushort v{{[0-9]+}}, v[{{[0-9:]+}}]
+; SRAMECC-ANY: s_setpc_b64
+
+; Unsupported target (gfx900) ignores sramecc flag and uses _d16_hi
+; GFX900-LABEL: {{^}}load_d16:
+; GFX900: s_waitcnt
+; GFX900: flat_load_short_d16_hi v0, v[{{[0-9:]+}}]
+; GFX900: s_setpc_b64
+
+;--- on.ll
+define <2 x i16> @load_d16(<2 x i16> %vec, ptr %ptr) {
+ %val = load i16, ptr %ptr
+ %result = insertelement <2 x i16> %vec, i16 %val, i32 1
+ ret <2 x i16> %result
+}
+!llvm.module.flags = !{!0}
+!0 = !{i32 1, !"amdgpu.sramecc", i32 1}
+
+;--- off.ll
+define <2 x i16> @load_d16(<2 x i16> %vec, ptr %ptr) {
+ %val = load i16, ptr %ptr
+ %result = insertelement <2 x i16> %vec, i16 %val, i32 1
+ ret <2 x i16> %result
+}
+!llvm.module.flags = !{!0}
+!0 = !{i32 1, !"amdgpu.sramecc", i32 0}
+
+;--- absent.ll
+define <2 x i16> @load_d16(<2 x i16> %vec, ptr %ptr) {
+ %val = load i16, ptr %ptr
+ %result = insertelement <2 x i16> %vec, i16 %val, i32 1
+ ret <2 x i16> %result
+}
diff --git a/llvm/test/CodeGen/AMDGPU/module-flag-xnack-sramecc-combined.ll b/llvm/test/CodeGen/AMDGPU/module-flag-xnack-sramecc-combined.ll
new file mode 100644
index 0000000000000..3d406834980e5
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/module-flag-xnack-sramecc-combined.ll
@@ -0,0 +1,15 @@
+; Test that xnack and sramecc target ID come from module flags
+;
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a < %s | FileCheck %s
+
+; Verify the target ID uses module flags (xnack+:sramecc-)
+; CHECK: .amdgcn_target "amdgcn-amd-amdhsa--gfx90a:sramecc-:xnack+"
+; CHECK: amdhsa.target: 'amdgcn-amd-amdhsa--gfx90a:sramecc-:xnack+'
+
+define void @foo() {
+ ret void
+}
+
+!llvm.module.flags = !{!0, !1}
+!0 = !{i32 1, !"amdgpu.xnack", i32 1}
+!1 = !{i32 1, !"amdgpu.sramecc", i32 0}
diff --git a/llvm/test/CodeGen/AMDGPU/module-flag-xnack.ll b/llvm/test/CodeGen/AMDGPU/module-flag-xnack.ll
new file mode 100644
index 0000000000000..ccce5da39a8f7
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/module-flag-xnack.ll
@@ -0,0 +1,75 @@
+; Test that .amdgcn_target directive includes xnack modifier based on module flag
+; Tests xnack+ (on), xnack- (off), and absent (Any) cases
+; Also tests that unsupported targets ignore the xnack module flag
+
+; RUN: split-file %s %t
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 < %t/on.ll | FileCheck --check-prefix=XNACK-ON %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 < %t/off.ll | FileCheck --check-prefix=XNACK-OFF %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 < %t/absent.ll | FileCheck --check-prefix=XNACK-ANY %s
+
+; Test that xnack module flag is ignored on targets that don't support it. gfx801 supports xnack, gfx803 does not.
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx801 < %t/on.ll | FileCheck --check-prefixes=CHECK,GFX801 %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx803 < %t/on.ll | FileCheck --check-prefixes=CHECK,GFX803 %s
+
+; Target directives for xnack supported target
+; XNACK-ON: .amdgcn_target "amdgcn-amd-amdhsa--gfx900:xnack+"
+; XNACK-OFF: .amdgcn_target "amdgcn-amd-amdhsa--gfx900:xnack-"
+; XNACK-ANY: .amdgcn_target "amdgcn-amd-amdhsa--gfx900"
+
+; GFX801: .amdgcn_target "amdgcn-amd-amdhsa--gfx801:xnack+"
+; GFX803: .amdgcn_target "amdgcn-amd-amdhsa--gfx803"
+
+; Check codegen impact - xnack affects register allocation
+; When xnack is on, first load must not overwrite the pointer argument
+; XNACK-ON-LABEL: {{^}}simple_clause:
+; XNACK-ON: flat_load_dword v4, v[0:1]
+; XNACK-ON-NEXT: flat_load_dword v5, v[2:3]
+
+; When xnack is off, first load can overwrite the pointer argument
+; XNACK-OFF-LABEL: {{^}}simple_clause:
+; XNACK-OFF: flat_load_dword v0, v[0:1]
+; XNACK-OFF-NEXT: flat_load_dword v1, v[2:3]
+
+; When xnack is not specified (Any), behavior is conservative (like on)
+; XNACK-ANY-LABEL: {{^}}simple_clause:
+; XNACK-ANY: flat_load_dword v4, v[0:1]
+; XNACK-ANY-NEXT: flat_load_dword v5, v[2:3]
+
+; Codegen for supported vs unsupported targets
+; CHECK-LABEL: {{^}}simple_clause:
+
+; First load must not overwrite the pointer argument on gfx801 (xnack supported)
+; GFX801: flat_load_dword v4, v[0:1]
+; GFX801-NEXT: flat_load_dword v5, v[2:3]
+
+; First load overwrites the pointer argument on gfx803 (xnack not supported)
+; GFX803: flat_load_dword v0, v[0:1]
+; GFX803-NEXT: flat_load_dword v1, v[2:3]
+
+;--- on.ll
+define i32 @simple_clause(ptr %ptr0, ptr %ptr1) {
+ %val0 = load i32, ptr %ptr0
+ %val1 = load i32, ptr %ptr1
+ %add = add i32 %val0, %val1
+ ret i32 %add
+}
+!llvm.module.flags = !{!0}
+!0 = !{i32 1, !"amdgpu.xnack", i32 1}
+
+;--- off.ll
+define i32 @simple_clause(ptr %ptr0, ptr %ptr1) {
+ %val0 = load i32, ptr %ptr0
+ %val1 = load i32, ptr %ptr1
+ %add = add i32 %val0, %val1
+ ret i32 %add
+}
+!llvm.module.flags = !{!0}
+!0 = !{i32 1, !"amdgpu.xnack", i32 0}
+
+;--- absent.ll
+define i32 @simple_clause(ptr %ptr0, ptr %ptr1) {
+ %val0 = load i32, ptr %ptr0
+ %val1 = load i32, ptr %ptr1
+ %add = add i32 %val0, %val1
+ ret i32 %add
+}
diff --git a/llvm/test/CodeGen/AMDGPU/nsa-reassign.ll b/llvm/test/CodeGen/AMDGPU/nsa-reassign.ll
index 4546d6c28286e..ca83283b67f07 100644
--- a/llvm/test/CodeGen/AMDGPU/nsa-reassign.ll
+++ b/llvm/test/CodeGen/AMDGPU/nsa-reassign.ll
@@ -1,4 +1,4 @@
-; RUN: llc -mtriple=amdgcn -mcpu=gfx1010 -mattr=-xnack -enable-misched=0 < %s | FileCheck -check-prefix=GCN %s
+; RUN: llc -mtriple=amdgcn -mcpu=gfx1010 --amdgpu-xnack=false -enable-misched=0 < %s | FileCheck -check-prefix=GCN %s
; GCN-LABEL: {{^}}sample_contig_nsa:
; GCN-DAG: image_sample_c_l v{{[0-9]+}}, v[{{[0-9]+:[0-9]+}}],
diff --git a/llvm/test/CodeGen/AMDGPU/nsa-vmem-hazard.mir b/llvm/test/CodeGen/AMDGPU/nsa-vmem-hazard.mir
index a968d4a1acfcf..459700d6549de 100644
--- a/llvm/test/CodeGen/AMDGPU/nsa-vmem-hazard.mir
+++ b/llvm/test/CodeGen/AMDGPU/nsa-vmem-hazard.mir
@@ -1,4 +1,4 @@
-# RUN: llc -mtriple=amdgcn -mcpu=gfx1010 -mattr=-xnack -verify-machineinstrs -run-pass post-RA-hazard-rec -o - %s | FileCheck -check-prefix=GCN %s
+# RUN: llc -mtriple=amdgcn -mcpu=gfx1010 --amdgpu-xnack=false -verify-machineinstrs -run-pass post-RA-hazard-rec -o - %s | FileCheck -check-prefix=GCN %s
# GCN-LABEL: name: hazard_image_sample_d_buf_off6
# GCN: IMAGE_SAMPLE
diff --git a/llvm/test/CodeGen/AMDGPU/occupancy-levels.ll b/llvm/test/CodeGen/AMDGPU/occupancy-levels.ll
index 2ede4248508ed..6d2a74702a39c 100644
--- a/llvm/test/CodeGen/AMDGPU/occupancy-levels.ll
+++ b/llvm/test/CodeGen/AMDGPU/occupancy-levels.ll
@@ -1,12 +1,7 @@
-; RUN: llc -mtriple=amdgcn -mcpu=gfx900 -mattr=-xnack < %s | FileCheck --check-prefixes=GCN,GFX9 %s
-; RUN: llc -mtriple=amdgcn -mcpu=gfx950 -mattr=-xnack < %s | FileCheck --check-prefixes=GCN,GFX950 %s
-; The amdhsa OS implicitly enables the trap handler, which reserves 16 SGPRs per
-; wave. The reported occupancy of SGPR-limited kernels must account for that, so
-; the same kernels reach a lower occupancy than on the non-amdhsa runs above.
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -mattr=-xnack < %s | FileCheck --check-prefixes=GCN,GFX9TRAP %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 -mattr=-xnack < %s | FileCheck --check-prefixes=GCN,GFX950TRAP %s
-; RUN: llc -mtriple=amdgcn -mcpu=gfx1010 -mattr=-xnack < %s | FileCheck --check-prefixes=GCN,GFX10,GFX10W32,GFX1010,GFX1010W32 %s
-; RUN: llc -mtriple=amdgcn -mcpu=gfx1010 -mattr=-xnack -mattr=+wavefrontsize64 < %s | FileCheck --check-prefixes=GCN,GFX10,GFX10W64,GFX1010,GFX1010W64 %s
+; RUN: llc -mtriple=amdgcn -mcpu=gfx900 --amdgpu-xnack=false < %s | FileCheck --check-prefixes=GCN,GFX9 %s
+; RUN: llc -mtriple=amdgcn -mcpu=gfx950 --amdgpu-xnack=false < %s | FileCheck --check-prefixes=GCN,GFX950 %s
+; RUN: llc -mtriple=amdgcn -mcpu=gfx1010 --amdgpu-xnack=false < %s | FileCheck --check-prefixes=GCN,GFX10,GFX10W32,GFX1010,GFX1010W32 %s
+; RUN: llc -mtriple=amdgcn -mcpu=gfx1010 --amdgpu-xnack=false -mattr=+wavefrontsize64 < %s | FileCheck --check-prefixes=GCN,GFX10,GFX10W64,GFX1010,GFX1010W64 %s
; RUN: llc -mtriple=amdgcn -mcpu=gfx1030 < %s | FileCheck --check-prefixes=GCN,GFX10,GFX10W32,GFX1030,GFX1030W32 %s
; RUN: llc -mtriple=amdgcn -mcpu=gfx1030 -mattr=+wavefrontsize64 < %s | FileCheck --check-prefixes=GCN,GFX10,GFX10W64,GFX1030,GFX1030W64 %s
; RUN: llc -mtriple=amdgcn -mcpu=gfx1100 < %s | FileCheck --check-prefixes=GCN,GFX1100,GFX1100W32 %s
diff --git a/llvm/test/CodeGen/AMDGPU/post-ra-soft-clause-dbg-info.ll b/llvm/test/CodeGen/AMDGPU/post-ra-soft-clause-dbg-info.ll
index 7a290a322e9e2..e9c7b0121ec21 100644
--- a/llvm/test/CodeGen/AMDGPU/post-ra-soft-clause-dbg-info.ll
+++ b/llvm/test/CodeGen/AMDGPU/post-ra-soft-clause-dbg-info.ll
@@ -1,5 +1,5 @@
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
-; RUN: llc -mtriple=amdgcn -mcpu=gfx900 -mattr=+xnack -amdgpu-max-memory-clause=0 < %s | FileCheck -enable-var-scope -check-prefix=GCN %s
+; RUN: llc -mtriple=amdgcn -mcpu=gfx900 --amdgpu-xnack=true -amdgpu-max-memory-clause=0 < %s | FileCheck -enable-var-scope -check-prefix=GCN %s
; Test the behavior of the post-RA soft clause bundler in the presence
; of debug info. The debug info should not interfere with the
diff --git a/llvm/test/CodeGen/AMDGPU/s_addk_i32.ll b/llvm/test/CodeGen/AMDGPU/s_addk_i32.ll
index f14a5cc19774d..76eda75e81cc1 100644
--- a/llvm/test/CodeGen/AMDGPU/s_addk_i32.ll
+++ b/llvm/test/CodeGen/AMDGPU/s_addk_i32.ll
@@ -1,5 +1,5 @@
; RUN: llc -mtriple=amdgcn--amdpal -mcpu=tahiti < %s | FileCheck -check-prefix=SI %s
-; RUN: llc -mtriple=amdgcn--amdpal -mcpu=tonga -mattr=-flat-for-global,-xnack < %s | FileCheck -check-prefix=SI %s
+; RUN: llc -mtriple=amdgcn--amdpal -mcpu=tonga -mattr=-flat-for-global --amdgpu-xnack=false < %s | FileCheck -check-prefix=SI %s
; TODO: Some of those tests fail with OS == amdhsa due to unreasonable register
; allocation differences.
diff --git a/llvm/test/CodeGen/AMDGPU/s_mulk_i32.ll b/llvm/test/CodeGen/AMDGPU/s_mulk_i32.ll
index ab98e8125e801..dc2f60f8ce588 100644
--- a/llvm/test/CodeGen/AMDGPU/s_mulk_i32.ll
+++ b/llvm/test/CodeGen/AMDGPU/s_mulk_i32.ll
@@ -1,6 +1,6 @@
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
; RUN: llc -mtriple=amdgcn--amdpal -mcpu=tahiti < %s | FileCheck -check-prefix=GFX6 %s
-; RUN: llc -mtriple=amdgcn--amdpal -mcpu=tonga -mattr=-flat-for-global,-xnack < %s | FileCheck -check-prefix=GFX8 %s
+; RUN: llc -mtriple=amdgcn--amdpal -mcpu=tonga -mattr=-flat-for-global --amdgpu-xnack=false < %s | FileCheck -check-prefix=GFX8 %s
define amdgpu_kernel void @s_mulk_i32_k0(ptr addrspace(1) %out, i32 %b) {
; GFX6-LABEL: s_mulk_i32_k0:
diff --git a/llvm/test/CodeGen/AMDGPU/schedule-amdgpu-tracker-physreg-crash.ll b/llvm/test/CodeGen/AMDGPU/schedule-amdgpu-tracker-physreg-crash.ll
index f70cd6816a966..775a3b53ad711 100644
--- a/llvm/test/CodeGen/AMDGPU/schedule-amdgpu-tracker-physreg-crash.ll
+++ b/llvm/test/CodeGen/AMDGPU/schedule-amdgpu-tracker-physreg-crash.ll
@@ -1,5 +1,5 @@
-; RUN: not llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -mattr=+xnack -amdgpu-use-amdgpu-trackers=1 2>&1 < %s | FileCheck -check-prefixes=ERR-GCNTRACKERS %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -mattr=+xnack 2>&1 < %s | FileCheck -check-prefixes=GCN %s
+; RUN: not llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 --amdgpu-xnack=true -amdgpu-use-amdgpu-trackers=1 2>&1 < %s | FileCheck -check-prefixes=ERR-GCNTRACKERS %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 --amdgpu-xnack=true 2>&1 < %s | FileCheck -check-prefixes=GCN %s
%asm.output = type { <16 x i32>, <16 x i32>, <16 x i32>, <8 x i32>, <2 x i32>, i32, ; sgprs
<16 x i32>, <7 x i32>, ; vgprs
diff --git a/llvm/test/CodeGen/AMDGPU/soft-clause-dbg-value.mir b/llvm/test/CodeGen/AMDGPU/soft-clause-dbg-value.mir
index af9ff4bae8292..e9157c89e3b74 100644
--- a/llvm/test/CodeGen/AMDGPU/soft-clause-dbg-value.mir
+++ b/llvm/test/CodeGen/AMDGPU/soft-clause-dbg-value.mir
@@ -1,6 +1,6 @@
# NOTE: Assertions have been autogenerated by utils/update_mir_test_checks.py
-# RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 -mattr=+xnack -run-pass=si-form-memory-clauses -verify-machineinstrs -o - %s | FileCheck %s
-# RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 -mattr=+xnack -passes="si-form-memory-clauses" -o - %s | FileCheck %s
+# RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 --amdgpu-xnack=true -run-pass=si-form-memory-clauses -verify-machineinstrs -o - %s | FileCheck %s
+# RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 --amdgpu-xnack=true -passes="si-form-memory-clauses" -o - %s | FileCheck %s
# Make sure that debug instructions do not change the bundling, and
# the dbg_values which break the clause are inserted after the new
diff --git a/llvm/test/CodeGen/AMDGPU/spill-scavenge-offset.ll b/llvm/test/CodeGen/AMDGPU/spill-scavenge-offset.ll
index 10b65e26875bb..3dd2d2b0dbad1 100644
--- a/llvm/test/CodeGen/AMDGPU/spill-scavenge-offset.ll
+++ b/llvm/test/CodeGen/AMDGPU/spill-scavenge-offset.ll
@@ -1,7 +1,7 @@
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
; RUN: llc -mtriple=amdgcn -mcpu=verde -enable-misched=0 -post-RA-scheduler=0 -amdgpu-spill-sgpr-to-vgpr=0 < %s | FileCheck -check-prefixes=CHECK,GFX6 %s
; RUN: llc -sgpr-regalloc=basic -vgpr-regalloc=basic -mtriple=amdgcn -mcpu=tonga -enable-misched=0 -post-RA-scheduler=0 -amdgpu-spill-sgpr-to-vgpr=0 < %s | FileCheck --check-prefix=CHECK %s
-; RUN: llc -mtriple=amdgcn -mattr=-xnack,+enable-flat-scratch -mcpu=gfx900 -enable-misched=0 -post-RA-scheduler=0 -amdgpu-spill-sgpr-to-vgpr=0 < %s | FileCheck -check-prefixes=CHECK,GFX9-FLATSCR,FLATSCR %s
+; RUN: llc -mtriple=amdgcn --amdgpu-xnack=false -mattr=+enable-flat-scratch -mcpu=gfx900 -enable-misched=0 -post-RA-scheduler=0 -amdgpu-spill-sgpr-to-vgpr=0 < %s | FileCheck -check-prefixes=CHECK,GFX9-FLATSCR,FLATSCR %s
; RUN: llc -mtriple=amdgcn -mcpu=gfx1030 -enable-misched=0 -post-RA-scheduler=0 -amdgpu-spill-sgpr-to-vgpr=0 -mattr=+enable-flat-scratch < %s | FileCheck -check-prefixes=CHECK,GFX10-FLATSCR,FLATSCR %s
;
; There is something about Tonga that causes this test to spend a lot of time
diff --git a/llvm/test/CodeGen/AMDGPU/sram-ecc-default.ll b/llvm/test/CodeGen/AMDGPU/sram-ecc-default.ll
index 18a991c7bab4e..56d101dd2f003 100644
--- a/llvm/test/CodeGen/AMDGPU/sram-ecc-default.ll
+++ b/llvm/test/CodeGen/AMDGPU/sram-ecc-default.ll
@@ -1,10 +1,10 @@
+; Flag is ignored on targets without sramecc support (like gfx900)
; RUN: llc -mtriple=amdgcn -mcpu=gfx900 < %s | FileCheck -check-prefixes=GCN,NO-ECC %s
-; RUN: llc -mtriple=amdgcn -mcpu=gfx900 -mattr=+sramecc < %s | FileCheck -check-prefixes=GCN,NO-ECC %s
-; RUN: llc -mtriple=amdgcn -mcpu=gfx900 -mattr=-sramecc < %s | FileCheck -check-prefixes=GCN,NO-ECC %s
-; RUN: llc -mtriple=amdgcn -mcpu=gfx902 -mattr=+sramecc < %s | FileCheck -check-prefixes=GCN,NO-ECC %s
-; RUN: llc -mtriple=amdgcn -mcpu=gfx904 -mattr=+sramecc < %s | FileCheck -check-prefixes=GCN,NO-ECC %s
-; RUN: llc -mtriple=amdgcn -mcpu=gfx906 -mattr=+sramecc < %s | FileCheck -check-prefixes=GCN,ECC %s
-; RUN: llc -mtriple=amdgcn -mcpu=gfx906 -mattr=-sramecc < %s | FileCheck -check-prefixes=GCN,NO-ECC %s
+; RUN: llc -mtriple=amdgcn -mcpu=gfx900 -amdgpu-sramecc=1 < %s | FileCheck -check-prefixes=GCN,NO-ECC %s
+; RUN: llc -mtriple=amdgcn -mcpu=gfx900 -amdgpu-sramecc=0 < %s | FileCheck -check-prefixes=GCN,NO-ECC %s
+
+; RUN: llc -mtriple=amdgcn -mcpu=gfx906 -amdgpu-sramecc=1 < %s | FileCheck -check-prefixes=GCN,ECC %s
+; RUN: llc -mtriple=amdgcn -mcpu=gfx906 -amdgpu-sramecc=0 < %s | FileCheck -check-prefixes=GCN,NO-ECC %s
; RUN: llc -mtriple=amdgcn -mcpu=gfx1250 < %s | FileCheck -check-prefixes=GCN,ECC %s
; Make sure the correct set of targets are marked with
diff --git a/llvm/test/CodeGen/AMDGPU/sramecc-subtarget-feature-disabled.ll b/llvm/test/CodeGen/AMDGPU/sramecc-subtarget-feature-disabled.ll
index 65b289bcd29d9..444e70f63c15d 100644
--- a/llvm/test/CodeGen/AMDGPU/sramecc-subtarget-feature-disabled.ll
+++ b/llvm/test/CodeGen/AMDGPU/sramecc-subtarget-feature-disabled.ll
@@ -1,10 +1,8 @@
-; RUN: llc -mtriple=amdgcn -mcpu=gfx700 -debug-only=gcn-subtarget -o - %s 2>&1 | FileCheck --check-prefix=WARN %s
; RUN: llc -mtriple=amdgcn -mcpu=gfx906 -debug-only=gcn-subtarget -o - %s 2>&1 | FileCheck --check-prefix=OFF %s
; RUN: llc -mtriple=amdgcn -mcpu=gfx908 -debug-only=gcn-subtarget -o - %s 2>&1 | FileCheck --check-prefix=OFF %s
; REQUIRES: asserts
-; WARN: warning: sramecc 'Off' was requested for a processor that does not support it!
; OFF: sramecc setting for subtarget: Off
define void @sramecc-subtarget-feature-disabled() #0 {
@@ -12,3 +10,6 @@ define void @sramecc-subtarget-feature-disabled() #0 {
}
attributes #0 = { "target-features"="-sramecc" }
+
+!llvm.module.flags = !{!0}
+!0 = !{i32 1, !"amdgpu.sramecc", i32 0}
diff --git a/llvm/test/CodeGen/AMDGPU/sramecc-subtarget-feature-enabled.ll b/llvm/test/CodeGen/AMDGPU/sramecc-subtarget-feature-enabled.ll
index ea6e456c55922..db2b3208873d3 100644
--- a/llvm/test/CodeGen/AMDGPU/sramecc-subtarget-feature-enabled.ll
+++ b/llvm/test/CodeGen/AMDGPU/sramecc-subtarget-feature-enabled.ll
@@ -1,14 +1,15 @@
-; RUN: llc -mtriple=amdgcn -mcpu=gfx700 -debug-only=gcn-subtarget -o - %s 2>&1 | FileCheck --check-prefix=WARN %s
; RUN: llc -mtriple=amdgcn -mcpu=gfx906 -debug-only=gcn-subtarget -o - %s 2>&1 | FileCheck --check-prefix=ON %s
; RUN: llc -mtriple=amdgcn -mcpu=gfx908 -debug-only=gcn-subtarget -o - %s 2>&1 | FileCheck --check-prefix=ON %s
; RUN: llc -mtriple=amdgcn -mcpu=gfx1250 -debug-only=gcn-subtarget -o - %s 2>&1 | FileCheck --check-prefix=ON %s
; REQUIRES: asserts
-; WARN: warning: sramecc 'On' was requested for a processor that does not support it!
; ON: sramecc setting for subtarget: On
define void @sramecc-subtarget-feature-enabled() #0 {
ret void
}
attributes #0 = { "target-features"="+sramecc" }
+
+!llvm.module.flags = !{!0}
+!0 = !{i32 1, !"amdgpu.sramecc", i32 1}
diff --git a/llvm/test/CodeGen/AMDGPU/target-id-xnack-always-on.ll b/llvm/test/CodeGen/AMDGPU/target-id-xnack-always-on.ll
index 13d13c875b8aa..96034a2ed7280 100644
--- a/llvm/test/CodeGen/AMDGPU/target-id-xnack-always-on.ll
+++ b/llvm/test/CodeGen/AMDGPU/target-id-xnack-always-on.ll
@@ -6,13 +6,13 @@
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1251 < %s | FileCheck --check-prefix=CHECK %s
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx12-5-generic < %s | FileCheck --check-prefix=CHECK %s
-; Even with -mattr=+xnack or -mattr=-xnack, the target ID doesn't change
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 -mattr=+xnack < %s | FileCheck %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 -mattr=-xnack < %s | FileCheck %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1251 -mattr=+xnack < %s | FileCheck %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1251 -mattr=-xnack < %s | FileCheck %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx12-5-generic -mattr=+xnack < %s | FileCheck %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx12-5-generic -mattr=-xnack < %s | FileCheck %s
+; Even with --amdgpu-xnack=true or --amdgpu-xnack=false, the target ID doesn't change
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 --amdgpu-xnack=true < %s | FileCheck %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 --amdgpu-xnack=false < %s | FileCheck %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1251 --amdgpu-xnack=true < %s | FileCheck %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1251 --amdgpu-xnack=false < %s | FileCheck %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx12-5-generic --amdgpu-xnack=true < %s | FileCheck %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx12-5-generic --amdgpu-xnack=false < %s | FileCheck %s
; CHECK: .amdgcn_target "amdgcn-amd-amdhsa--gfx{{1250|1251|12-5-generic}}"
diff --git a/llvm/test/CodeGen/AMDGPU/tid-kd-xnack-off.ll b/llvm/test/CodeGen/AMDGPU/tid-kd-xnack-off.ll
index 2922424704edc..0474f7727d2ce 100644
--- a/llvm/test/CodeGen/AMDGPU/tid-kd-xnack-off.ll
+++ b/llvm/test/CodeGen/AMDGPU/tid-kd-xnack-off.ll
@@ -1,6 +1,6 @@
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a -mattr=-xnack < %s | FileCheck --check-prefixes=ASM %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a -mattr=-xnack --filetype=obj < %s | llvm-objdump -s -j .rodata - | FileCheck --check-prefixes=OBJ %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a -mattr=-xnack --filetype=obj < %s | llvm-readelf --notes - | FileCheck --check-prefixes=ELF %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a < %s | FileCheck --check-prefixes=ASM %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a --filetype=obj < %s | llvm-objdump -s -j .rodata - | FileCheck --check-prefixes=OBJ %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a --filetype=obj < %s | llvm-readelf --notes - | FileCheck --check-prefixes=ELF %s
; TODO: Update to check for granulated sgpr count directive once one is added.
@@ -25,5 +25,6 @@ entry:
attributes #0 = { "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-cluster-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-cluster-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-cluster-id-z" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }
-!llvm.module.flags = !{!0}
+!llvm.module.flags = !{!0, !1}
!0 = !{i32 1, !"amdhsa_code_object_version", i32 400}
+!1 = !{i32 1, !"amdgpu.xnack", i32 0}
diff --git a/llvm/test/CodeGen/AMDGPU/tid-kd-xnack-on.ll b/llvm/test/CodeGen/AMDGPU/tid-kd-xnack-on.ll
index aedb5f9106ec8..b1785073ba2a5 100644
--- a/llvm/test/CodeGen/AMDGPU/tid-kd-xnack-on.ll
+++ b/llvm/test/CodeGen/AMDGPU/tid-kd-xnack-on.ll
@@ -1,6 +1,6 @@
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a -mattr=+xnack < %s | FileCheck --check-prefixes=ASM %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a -mattr=+xnack --filetype=obj < %s | llvm-objdump -s -j .rodata - | FileCheck --check-prefixes=OBJ %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a -mattr=+xnack --filetype=obj < %s | llvm-readelf --notes - | FileCheck --check-prefixes=ELF %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a --amdgpu-xnack=true < %s | FileCheck --check-prefixes=ASM %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a --amdgpu-xnack=true --filetype=obj < %s | llvm-objdump -s -j .rodata - | FileCheck --check-prefixes=OBJ %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a --amdgpu-xnack=true --filetype=obj < %s | llvm-readelf --notes - | FileCheck --check-prefixes=ELF %s
; TODO: Update to check for granulated sgpr count directive once one is added.
diff --git a/llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-all-off.ll b/llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-all-off.ll
index 08dd90250d0b4..69cd6bdc1f75f 100644
--- a/llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-all-off.ll
+++ b/llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-all-off.ll
@@ -42,5 +42,6 @@ entry:
}
attributes #0 = { "target-features"="-xnack" }
-!llvm.module.flags = !{!0}
+!llvm.module.flags = !{!0, !1}
!0 = !{i32 1, !"amdhsa_code_object_version", i32 CODE_OBJECT_VERSION}
+!1 = !{i32 1, !"amdgpu.xnack", i32 0}
diff --git a/llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-all-on.ll b/llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-all-on.ll
index a8340ddadaaf7..d6f34c3108df8 100644
--- a/llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-all-on.ll
+++ b/llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-all-on.ll
@@ -43,5 +43,6 @@ entry:
attributes #0 = { "target-features"="+xnack" }
-!llvm.module.flags = !{!0}
+!llvm.module.flags = !{!0, !1}
!0 = !{i32 1, !"amdhsa_code_object_version", i32 CODE_OBJECT_VERSION}
+!1 = !{i32 1, !"amdgpu.xnack", i32 1}
diff --git a/llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-any-off-1.ll b/llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-any-off-1.ll
index aefcfac23ff5d..708012974bcfb 100644
--- a/llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-any-off-1.ll
+++ b/llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-any-off-1.ll
@@ -43,5 +43,6 @@ entry:
attributes #0 = { "target-features"="-xnack" }
-!llvm.module.flags = !{!0}
+!llvm.module.flags = !{!0, !1}
!0 = !{i32 1, !"amdhsa_code_object_version", i32 CODE_OBJECT_VERSION}
+!1 = !{i32 1, !"amdgpu.xnack", i32 0}
diff --git a/llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-any-off-2.ll b/llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-any-off-2.ll
index 6005c31622405..3438116ffefd7 100644
--- a/llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-any-off-2.ll
+++ b/llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-any-off-2.ll
@@ -43,5 +43,6 @@ entry:
attributes #0 = { "target-features"="-xnack" }
-!llvm.module.flags = !{!0}
+!llvm.module.flags = !{!0, !1}
!0 = !{i32 1, !"amdhsa_code_object_version", i32 CODE_OBJECT_VERSION}
+!1 = !{i32 1, !"amdgpu.xnack", i32 0}
diff --git a/llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-any-on-1.ll b/llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-any-on-1.ll
index 328f56fb841b8..9a232d5c3cec0 100644
--- a/llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-any-on-1.ll
+++ b/llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-any-on-1.ll
@@ -43,5 +43,6 @@ entry:
attributes #0 = { "target-features"="+xnack" }
-!llvm.module.flags = !{!0}
+!llvm.module.flags = !{!0, !1}
!0 = !{i32 1, !"amdhsa_code_object_version", i32 CODE_OBJECT_VERSION}
+!1 = !{i32 1, !"amdgpu.xnack", i32 1}
diff --git a/llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-any-on-2.ll b/llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-any-on-2.ll
index c50dd8b2fec7a..2c41517c57c0d 100644
--- a/llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-any-on-2.ll
+++ b/llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-any-on-2.ll
@@ -43,5 +43,6 @@ entry:
attributes #0 = { "target-features"="+xnack" }
-!llvm.module.flags = !{!0}
+!llvm.module.flags = !{!0, !1}
!0 = !{i32 1, !"amdhsa_code_object_version", i32 CODE_OBJECT_VERSION}
+!1 = !{i32 1, !"amdgpu.xnack", i32 1}
diff --git a/llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-invalid-any-off-on.ll b/llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-invalid-any-off-on.ll
deleted file mode 100644
index 0f54d783484dd..0000000000000
--- a/llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-invalid-any-off-on.ll
+++ /dev/null
@@ -1,24 +0,0 @@
-; RUN: not llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 < %s 2>&1 | FileCheck --check-prefixes=ERR %s
-
-; ERR: error: xnack setting of 'func2' function does not match module xnack setting
-
-define void @func0() {
-entry:
- ret void
-}
-
-define void @func1() #0 {
-entry:
- ret void
-}
-
-define void @func2() #1 {
-entry:
- ret void
-}
-
-attributes #0 = { "target-features"="-xnack" }
-attributes #1 = { "target-features"="+xnack" }
-
-!llvm.module.flags = !{!0}
-!0 = !{i32 1, !"amdhsa_code_object_version", i32 400}
diff --git a/llvm/test/CodeGen/AMDGPU/tid-one-func-xnack-off.ll b/llvm/test/CodeGen/AMDGPU/tid-one-func-xnack-off.ll
index 60ff8b2dbb5eb..cbd870c32b4dd 100644
--- a/llvm/test/CodeGen/AMDGPU/tid-one-func-xnack-off.ll
+++ b/llvm/test/CodeGen/AMDGPU/tid-one-func-xnack-off.ll
@@ -33,5 +33,6 @@ entry:
attributes #0 = { "target-features"="-xnack" }
-!llvm.module.flags = !{!0}
+!llvm.module.flags = !{!0, !1}
!0 = !{i32 1, !"amdhsa_code_object_version", i32 CODE_OBJECT_VERSION}
+!1 = !{i32 1, !"amdgpu.xnack", i32 0}
diff --git a/llvm/test/CodeGen/AMDGPU/tid-one-func-xnack-on.ll b/llvm/test/CodeGen/AMDGPU/tid-one-func-xnack-on.ll
index e04629a24209e..8713cc78e6d4d 100644
--- a/llvm/test/CodeGen/AMDGPU/tid-one-func-xnack-on.ll
+++ b/llvm/test/CodeGen/AMDGPU/tid-one-func-xnack-on.ll
@@ -33,5 +33,6 @@ entry:
attributes #0 = { "target-features"="+xnack" }
-!llvm.module.flags = !{!0}
+!llvm.module.flags = !{!0, !1}
!0 = !{i32 1, !"amdhsa_code_object_version", i32 CODE_OBJECT_VERSION}
+!1 = !{i32 1, !"amdgpu.xnack", i32 1}
diff --git a/llvm/test/CodeGen/AMDGPU/xnack-subtarget-feature-disabled.ll b/llvm/test/CodeGen/AMDGPU/xnack-subtarget-feature-disabled.ll
index 111feb4503b20..8d281e66ae661 100644
--- a/llvm/test/CodeGen/AMDGPU/xnack-subtarget-feature-disabled.ll
+++ b/llvm/test/CodeGen/AMDGPU/xnack-subtarget-feature-disabled.ll
@@ -1,14 +1,10 @@
-; RUN: llc -mtriple=amdgcn -mcpu=gfx600 -debug-only=gcn-subtarget -filetype=null %s 2>&1 | FileCheck --check-prefix=WARN %s
-; RUN: llc -mtriple=amdgcn -mcpu=gfx700 -debug-only=gcn-subtarget -filetype=null %s 2>&1 | FileCheck --check-prefix=WARN %s
; RUN: llc -mtriple=amdgcn -mcpu=gfx801 -debug-only=gcn-subtarget -o - %s 2>&1 | FileCheck --check-prefix=OFF %s
; RUN: llc -mtriple=amdgcn -mcpu=gfx900 -debug-only=gcn-subtarget -o - %s 2>&1 | FileCheck --check-prefix=OFF %s
; RUN: llc -mtriple=amdgcn -mcpu=gfx906 -debug-only=gcn-subtarget -o - %s 2>&1 | FileCheck --check-prefix=OFF %s
; RUN: llc -mtriple=amdgcn -mcpu=gfx1010 -debug-only=gcn-subtarget -o - %s 2>&1 | FileCheck --check-prefix=OFF %s
-; RUN: llc -mtriple=amdgcn -mcpu=gfx1100 -debug-only=gcn-subtarget -o - %s 2>&1 | FileCheck --check-prefix=WARN %s
; REQUIRES: asserts
-; WARN: warning: xnack 'Off' was requested for a processor that does not support it!
; OFF: xnack setting for subtarget: Off
define void @xnack-subtarget-feature-disabled() #0 {
@@ -16,3 +12,6 @@ define void @xnack-subtarget-feature-disabled() #0 {
}
attributes #0 = { "target-features"="-xnack" }
+
+!llvm.module.flags = !{!0}
+!0 = !{i32 1, !"amdgpu.xnack", i32 0}
diff --git a/llvm/test/CodeGen/AMDGPU/xnack-subtarget-feature-enabled.ll b/llvm/test/CodeGen/AMDGPU/xnack-subtarget-feature-enabled.ll
index 3a1555fa367cf..96f52cd902237 100644
--- a/llvm/test/CodeGen/AMDGPU/xnack-subtarget-feature-enabled.ll
+++ b/llvm/test/CodeGen/AMDGPU/xnack-subtarget-feature-enabled.ll
@@ -1,17 +1,16 @@
-; RUN: llc -mtriple=amdgcn -mcpu=gfx600 -debug-only=gcn-subtarget -filetype=null %s 2>&1 | FileCheck --check-prefix=WARN %s
-; RUN: llc -mtriple=amdgcn -mcpu=gfx700 -debug-only=gcn-subtarget -filetype=null %s 2>&1 | FileCheck --check-prefix=WARN %s
; RUN: llc -mtriple=amdgcn -mcpu=gfx801 -debug-only=gcn-subtarget -o - %s 2>&1 | FileCheck --check-prefix=ON %s
; RUN: llc -mtriple=amdgcn -mcpu=gfx900 -debug-only=gcn-subtarget -o - %s 2>&1 | FileCheck --check-prefix=ON %s
; RUN: llc -mtriple=amdgcn -mcpu=gfx906 -debug-only=gcn-subtarget -o - %s 2>&1 | FileCheck --check-prefix=ON %s
; RUN: llc -mtriple=amdgcn -mcpu=gfx1010 -debug-only=gcn-subtarget -o - %s 2>&1 | FileCheck --check-prefix=ON %s
-; RUN: llc -mtriple=amdgcn -mcpu=gfx1100 -debug-only=gcn-subtarget -o - %s 2>&1 | FileCheck --check-prefix=WARN %s
; REQUIRES: asserts
-; WARN: warning: xnack 'On' was requested for a processor that does not support it!
; ON: xnack setting for subtarget: On
define void @xnack-subtarget-feature-enabled() #0 {
ret void
}
attributes #0 = { "target-features"="+xnack" }
+
+!llvm.module.flags = !{!0}
+!0 = !{i32 1, !"amdgpu.xnack", i32 1}
diff --git a/llvm/test/Linker/Inputs/amdgpu-sramecc-module-flag-0.ll b/llvm/test/Linker/Inputs/amdgpu-sramecc-module-flag-0.ll
new file mode 100644
index 0000000000000..d076baa83ce40
--- /dev/null
+++ b/llvm/test/Linker/Inputs/amdgpu-sramecc-module-flag-0.ll
@@ -0,0 +1,6 @@
+define void @input_off() {
+ ret void
+}
+
+!llvm.module.flags = !{!0}
+!0 = !{i32 1, !"amdgpu.sramecc", i32 0}
diff --git a/llvm/test/Linker/Inputs/amdgpu-sramecc-module-flag-1.ll b/llvm/test/Linker/Inputs/amdgpu-sramecc-module-flag-1.ll
new file mode 100644
index 0000000000000..38915c3cd82eb
--- /dev/null
+++ b/llvm/test/Linker/Inputs/amdgpu-sramecc-module-flag-1.ll
@@ -0,0 +1,6 @@
+define void @input_on() {
+ ret void
+}
+
+!llvm.module.flags = !{!0}
+!0 = !{i32 1, !"amdgpu.sramecc", i32 1}
diff --git a/llvm/test/Linker/Inputs/amdgpu-sramecc-module-flag-any.ll b/llvm/test/Linker/Inputs/amdgpu-sramecc-module-flag-any.ll
new file mode 100644
index 0000000000000..3e4b92c69b389
--- /dev/null
+++ b/llvm/test/Linker/Inputs/amdgpu-sramecc-module-flag-any.ll
@@ -0,0 +1,6 @@
+define void @input_any() {
+ ret void
+}
+
+!llvm.module.flags = !{!0}
+!0 = !{i32 7, !"PIC Level", i32 2}
diff --git a/llvm/test/Linker/Inputs/amdgpu-xnack-module-flag-0.ll b/llvm/test/Linker/Inputs/amdgpu-xnack-module-flag-0.ll
new file mode 100644
index 0000000000000..f90037f736578
--- /dev/null
+++ b/llvm/test/Linker/Inputs/amdgpu-xnack-module-flag-0.ll
@@ -0,0 +1,6 @@
+define void @input_off() {
+ ret void
+}
+
+!llvm.module.flags = !{!0}
+!0 = !{i32 1, !"amdgpu.xnack", i32 0}
diff --git a/llvm/test/Linker/Inputs/amdgpu-xnack-module-flag-1.ll b/llvm/test/Linker/Inputs/amdgpu-xnack-module-flag-1.ll
new file mode 100644
index 0000000000000..c2b16ff433170
--- /dev/null
+++ b/llvm/test/Linker/Inputs/amdgpu-xnack-module-flag-1.ll
@@ -0,0 +1,6 @@
+define void @input_on() {
+ ret void
+}
+
+!llvm.module.flags = !{!0}
+!0 = !{i32 1, !"amdgpu.xnack", i32 1}
diff --git a/llvm/test/Linker/Inputs/amdgpu-xnack-module-flag-any.ll b/llvm/test/Linker/Inputs/amdgpu-xnack-module-flag-any.ll
new file mode 100644
index 0000000000000..3e4b92c69b389
--- /dev/null
+++ b/llvm/test/Linker/Inputs/amdgpu-xnack-module-flag-any.ll
@@ -0,0 +1,6 @@
+define void @input_any() {
+ ret void
+}
+
+!llvm.module.flags = !{!0}
+!0 = !{i32 7, !"PIC Level", i32 2}
diff --git a/llvm/test/Linker/amdgpu-sramecc-module-flag-0.ll b/llvm/test/Linker/amdgpu-sramecc-module-flag-0.ll
new file mode 100644
index 0000000000000..95d2680daf31b
--- /dev/null
+++ b/llvm/test/Linker/amdgpu-sramecc-module-flag-0.ll
@@ -0,0 +1,24 @@
+; Test that sramecc module flags are linked correctly with Module::Error behavior
+
+; RUN: llvm-link -S %s %S/Inputs/amdgpu-sramecc-module-flag-0.ll -o - | FileCheck --check-prefix=BOTH-OFF %s
+; RUN: not llvm-link -S %s %S/Inputs/amdgpu-sramecc-module-flag-1.ll -o /dev/null 2>&1 | FileCheck --check-prefix=CONFLICT %s
+; RUN: llvm-link -S %s %S/Inputs/amdgpu-sramecc-module-flag-any.ll -o - | FileCheck --check-prefix=ONE-OFF %s
+
+; Test disabled + disabled = disabled
+; BOTH-OFF: !llvm.module.flags = !{!0}
+; BOTH-OFF: !0 = !{i32 1, !"amdgpu.sramecc", i32 0}
+
+; Test disabled + enabled = error
+; CONFLICT: linking module flags 'amdgpu.sramecc': IDs have conflicting values
+
+; Test disabled + any = disabled
+; ONE-OFF: !llvm.module.flags = !{!0, !1}
+; ONE-OFF-DAG: !{{[0-9]}} = !{i32 1, !"amdgpu.sramecc", i32 0}
+; ONE-OFF-DAG: !{{[0-9]}} = !{i32 {{[0-9]+}}, !"PIC Level", i32 {{[0-9]+}}}
+
+define void @foo() {
+ ret void
+}
+
+!llvm.module.flags = !{!0}
+!0 = !{i32 1, !"amdgpu.sramecc", i32 0}
diff --git a/llvm/test/Linker/amdgpu-sramecc-module-flag-1.ll b/llvm/test/Linker/amdgpu-sramecc-module-flag-1.ll
new file mode 100644
index 0000000000000..c19c2d80c95ee
--- /dev/null
+++ b/llvm/test/Linker/amdgpu-sramecc-module-flag-1.ll
@@ -0,0 +1,23 @@
+; Test sramecc module flag linking with enabled flag
+; RUN: not llvm-link -S %s %S/Inputs/amdgpu-sramecc-module-flag-0.ll -o /dev/null 2>&1 | FileCheck --check-prefix=CONFLICT %s
+; RUN: llvm-link -S %s %S/Inputs/amdgpu-sramecc-module-flag-1.ll -o - | FileCheck --check-prefix=BOTH-ON %s
+; RUN: llvm-link -S %s %S/Inputs/amdgpu-sramecc-module-flag-any.ll -o - | FileCheck --check-prefix=ONE-ON %s
+
+; Test enabled + disabled = error
+; CONFLICT: linking module flags 'amdgpu.sramecc': IDs have conflicting values
+
+; Test enabled + enabled = enabled
+; BOTH-ON: !llvm.module.flags = !{!0}
+; BOTH-ON: !0 = !{i32 1, !"amdgpu.sramecc", i32 1}
+
+; Test enabled + any = enabled
+; ONE-ON: !llvm.module.flags = !{!0, !1}
+; ONE-ON-DAG: !{{[0-9]}} = !{i32 1, !"amdgpu.sramecc", i32 1}
+; ONE-ON-DAG: !{{[0-9]}} = !{i32 {{[0-9]+}}, !"PIC Level", i32 {{[0-9]+}}}
+
+define void @bar() {
+ ret void
+}
+
+!llvm.module.flags = !{!0}
+!0 = !{i32 1, !"amdgpu.sramecc", i32 1}
diff --git a/llvm/test/Linker/amdgpu-sramecc-module-flag-any.ll b/llvm/test/Linker/amdgpu-sramecc-module-flag-any.ll
new file mode 100644
index 0000000000000..29b158bd9c340
--- /dev/null
+++ b/llvm/test/Linker/amdgpu-sramecc-module-flag-any.ll
@@ -0,0 +1,25 @@
+; Test sramecc module flag linking with no flag (any)
+; RUN: llvm-link -S %s %S/Inputs/amdgpu-sramecc-module-flag-0.ll -o - | FileCheck --check-prefix=OTHER-OFF %s
+; RUN: llvm-link -S %s %S/Inputs/amdgpu-sramecc-module-flag-1.ll -o - | FileCheck --check-prefix=OTHER-ON %s
+; RUN: llvm-link -S %s %S/Inputs/amdgpu-sramecc-module-flag-any.ll -o - | FileCheck --check-prefix=BOTH-ANY %s
+
+; Test any + disabled = disabled
+; OTHER-OFF: !llvm.module.flags = !{!0, !1}
+; OTHER-OFF-DAG: !{{[0-9]}} = !{i32 {{[0-9]+}}, !"PIC Level", i32 {{[0-9]+}}}
+; OTHER-OFF-DAG: !{{[0-9]}} = !{i32 1, !"amdgpu.sramecc", i32 0}
+
+; Test any + enabled = enabled
+; OTHER-ON: !llvm.module.flags = !{!0, !1}
+; OTHER-ON-DAG: !{{[0-9]}} = !{i32 {{[0-9]+}}, !"PIC Level", i32 {{[0-9]+}}}
+; OTHER-ON-DAG: !{{[0-9]}} = !{i32 1, !"amdgpu.sramecc", i32 1}
+
+; Test any + any = any (no flag)
+; BOTH-ANY: !llvm.module.flags = !{!0}
+; BOTH-ANY: !0 = !{i32 {{[0-9]+}}, !"PIC Level", i32 2}
+
+define void @baz() {
+ ret void
+}
+
+!llvm.module.flags = !{!0}
+!0 = !{i32 7, !"PIC Level", i32 2}
diff --git a/llvm/test/Linker/amdgpu-xnack-module-flag-0.ll b/llvm/test/Linker/amdgpu-xnack-module-flag-0.ll
new file mode 100644
index 0000000000000..bb06186fc96d2
--- /dev/null
+++ b/llvm/test/Linker/amdgpu-xnack-module-flag-0.ll
@@ -0,0 +1,24 @@
+; Test that xnack module flags are linked correctly with Module::Error behavior
+
+; RUN: llvm-link -S %s %S/Inputs/amdgpu-xnack-module-flag-0.ll -o - | FileCheck --check-prefix=BOTH-OFF %s
+; RUN: not llvm-link -S %s %S/Inputs/amdgpu-xnack-module-flag-1.ll -o /dev/null 2>&1 | FileCheck --check-prefix=CONFLICT %s
+; RUN: llvm-link -S %s %S/Inputs/amdgpu-xnack-module-flag-any.ll -o - | FileCheck --check-prefix=ONE-OFF %s
+
+; Test disabled + disabled = disabled
+; BOTH-OFF: !llvm.module.flags = !{!0}
+; BOTH-OFF: !0 = !{i32 1, !"amdgpu.xnack", i32 0}
+
+; Test disabled + enabled = error
+; CONFLICT: linking module flags 'amdgpu.xnack': IDs have conflicting values
+
+; Test disabled + any = disabled
+; ONE-OFF: !llvm.module.flags = !{!0, !1}
+; ONE-OFF-DAG: !{{[0-9]}} = !{i32 1, !"amdgpu.xnack", i32 0}
+; ONE-OFF-DAG: !{{[0-9]}} = !{i32 {{[0-9]+}}, !"PIC Level", i32 {{[0-9]+}}}
+
+define void @foo() {
+ ret void
+}
+
+!llvm.module.flags = !{!0}
+!0 = !{i32 1, !"amdgpu.xnack", i32 0}
diff --git a/llvm/test/Linker/amdgpu-xnack-module-flag-1.ll b/llvm/test/Linker/amdgpu-xnack-module-flag-1.ll
new file mode 100644
index 0000000000000..a3f11c258219c
--- /dev/null
+++ b/llvm/test/Linker/amdgpu-xnack-module-flag-1.ll
@@ -0,0 +1,23 @@
+; Test xnack module flag linking with enabled flag
+; RUN: not llvm-link -S %s %S/Inputs/amdgpu-xnack-module-flag-0.ll -o /dev/null 2>&1 | FileCheck --check-prefix=CONFLICT %s
+; RUN: llvm-link -S %s %S/Inputs/amdgpu-xnack-module-flag-1.ll -o - | FileCheck --check-prefix=BOTH-ON %s
+; RUN: llvm-link -S %s %S/Inputs/amdgpu-xnack-module-flag-any.ll -o - | FileCheck --check-prefix=ONE-ON %s
+
+; Test enabled + disabled = error
+; CONFLICT: linking module flags 'amdgpu.xnack': IDs have conflicting values
+
+; Test enabled + enabled = enabled
+; BOTH-ON: !llvm.module.flags = !{!0}
+; BOTH-ON: !0 = !{i32 1, !"amdgpu.xnack", i32 1}
+
+; Test enabled + any = enabled
+; ONE-ON: !llvm.module.flags = !{!0, !1}
+; ONE-ON-DAG: !{{[0-9]}} = !{i32 1, !"amdgpu.xnack", i32 1}
+; ONE-ON-DAG: !{{[0-9]}} = !{i32 {{[0-9]+}}, !"PIC Level", i32 {{[0-9]+}}}
+
+define void @bar() {
+ ret void
+}
+
+!llvm.module.flags = !{!0}
+!0 = !{i32 1, !"amdgpu.xnack", i32 1}
diff --git a/llvm/test/Linker/amdgpu-xnack-module-flag-any.ll b/llvm/test/Linker/amdgpu-xnack-module-flag-any.ll
new file mode 100644
index 0000000000000..53ce0aebfeac0
--- /dev/null
+++ b/llvm/test/Linker/amdgpu-xnack-module-flag-any.ll
@@ -0,0 +1,25 @@
+; Test xnack module flag linking with no flag (any)
+; RUN: llvm-link -S %s %S/Inputs/amdgpu-xnack-module-flag-0.ll -o - | FileCheck --check-prefix=OTHER-OFF %s
+; RUN: llvm-link -S %s %S/Inputs/amdgpu-xnack-module-flag-1.ll -o - | FileCheck --check-prefix=OTHER-ON %s
+; RUN: llvm-link -S %s %S/Inputs/amdgpu-xnack-module-flag-any.ll -o - | FileCheck --check-prefix=BOTH-ANY %s
+
+; Test any + disabled = disabled
+; OTHER-OFF: !llvm.module.flags = !{!0, !1}
+; OTHER-OFF-DAG: !{{[0-9]}} = !{i32 {{[0-9]+}}, !"PIC Level", i32 {{[0-9]+}}}
+; OTHER-OFF-DAG: !{{[0-9]}} = !{i32 1, !"amdgpu.xnack", i32 0}
+
+; Test any + enabled = enabled
+; OTHER-ON: !llvm.module.flags = !{!0, !1}
+; OTHER-ON-DAG: !{{[0-9]}} = !{i32 {{[0-9]+}}, !"PIC Level", i32 {{[0-9]+}}}
+; OTHER-ON-DAG: !{{[0-9]}} = !{i32 1, !"amdgpu.xnack", i32 1}
+
+; Test any + any = any (no flag)
+; BOTH-ANY: !llvm.module.flags = !{!0}
+; BOTH-ANY: !0 = !{i32 {{[0-9]+}}, !"PIC Level", i32 2}
+
+define void @baz() {
+ ret void
+}
+
+!llvm.module.flags = !{!0}
+!0 = !{i32 7, !"PIC Level", i32 2}
diff --git a/llvm/test/MC/AMDGPU/xnack-mask.s b/llvm/test/MC/AMDGPU/xnack-mask.s
index a473050685525..a5c26a06cd028 100644
--- a/llvm/test/MC/AMDGPU/xnack-mask.s
+++ b/llvm/test/MC/AMDGPU/xnack-mask.s
@@ -1,7 +1,7 @@
// RUN: not llvm-mc -triple=amdgcn -mcpu=tahiti %s -filetype=null 2>&1 | FileCheck -check-prefix=NOSICIVI10 --implicit-check-not=error: %s
// RUN: not llvm-mc -triple=amdgcn -mcpu=hawaii %s -filetype=null 2>&1 | FileCheck -check-prefix=NOSICIVI10 --implicit-check-not=error: %s
// RUN: not llvm-mc -triple=amdgcn -mcpu=tonga %s -filetype=null 2>&1 | FileCheck -check-prefix=NOSICIVI10 --implicit-check-not=error: %s
-// RUN: not llvm-mc -triple=amdgcn -mcpu=gfx1001 -mattr=-xnack %s -filetype=null 2>&1 | FileCheck -check-prefix=NOSICIVI10 --implicit-check-not=error: %s
+// RUN: not llvm-mc -triple=amdgcn -mcpu=gfx1010 -mattr=-xnack %s -filetype=null 2>&1 | FileCheck -check-prefix=NOSICIVI10 --implicit-check-not=error: %s
// RUN: not llvm-mc -triple=amdgcn -mcpu=stoney -mattr=+xnack %s -filetype=null 2>&1 | FileCheck -check-prefix=XNACKERR --implicit-check-not=error: %s
// RUN: not llvm-mc -triple=amdgcn -mcpu=stoney -mattr=+xnack -show-encoding %s | FileCheck -check-prefix=XNACK %s
diff --git a/llvm/test/Verifier/AMDGPU/module-flag-sramecc.ll b/llvm/test/Verifier/AMDGPU/module-flag-sramecc.ll
new file mode 100644
index 0000000000000..14308479c04bc
--- /dev/null
+++ b/llvm/test/Verifier/AMDGPU/module-flag-sramecc.ll
@@ -0,0 +1,46 @@
+; Tests for IR verifier enforcement of the "amdgpu.sramecc" module flag.
+; The flag must use Module::Error (i32 1) merge behavior, carry a constant
+; integer value, and be 0 or 1.
+
+; RUN: split-file %s %t
+
+; --- Negative: wrong merge behavior (Max=7 instead of Error=1) ---
+; RUN: not llvm-as %t/wrong-behavior.ll --disable-output 2>&1 \
+; RUN: | FileCheck %s --check-prefix=WRONG-BEHAVIOR
+
+; --- Negative: non-integer value ---
+; RUN: not llvm-as %t/non-integer.ll --disable-output 2>&1 \
+; RUN: | FileCheck %s --check-prefix=NON-INT
+
+; --- Negative: missing value ---
+; RUN: not llvm-as %t/missing-value.ll --disable-output 2>&1 \
+; RUN: | FileCheck %s --check-prefix=MISSING-VALUE
+
+; --- Negative: value out of range (2 is not 0 or 1) ---
+; RUN: not llvm-as %t/out-of-range.ll --disable-output 2>&1 \
+; RUN: | FileCheck %s --check-prefix=RANGE
+
+; WRONG-BEHAVIOR: 'amdgpu.sramecc' module flag must use 'error' merge behaviour
+; NON-INT: 'amdgpu.sramecc' module flag must have a constant integer value
+; MISSING-VALUE: incorrect number of operands in module flag
+; RANGE: 'amdgpu.sramecc' module flag must be 0 or 1
+
+;--- wrong-behavior.ll
+; Max (i32 7) is not Error (i32 1).
+!0 = !{i32 7, !"amdgpu.sramecc", i32 1}
+!llvm.module.flags = !{!0}
+
+;--- non-integer.ll
+; Error behavior but float value instead of integer.
+!0 = !{i32 1, !"amdgpu.sramecc", float 1.0}
+!llvm.module.flags = !{!0}
+
+;--- missing-value.ll
+; Missing value field.
+!0 = !{i32 1, !"amdgpu.sramecc"}
+!llvm.module.flags = !{!0}
+
+;--- out-of-range.ll
+; Value 2 is out of range (must be 0 or 1).
+!0 = !{i32 1, !"amdgpu.sramecc", i32 2}
+!llvm.module.flags = !{!0}
diff --git a/llvm/test/Verifier/AMDGPU/module-flag-xnack.ll b/llvm/test/Verifier/AMDGPU/module-flag-xnack.ll
new file mode 100644
index 0000000000000..23168ab5e3869
--- /dev/null
+++ b/llvm/test/Verifier/AMDGPU/module-flag-xnack.ll
@@ -0,0 +1,46 @@
+; Tests for IR verifier enforcement of the "amdgpu.xnack" module flag.
+; The flag must use Module::Error (i32 1) merge behavior, carry a constant
+; integer value, and be 0 or 1.
+
+; RUN: split-file %s %t
+
+; --- Negative: wrong merge behavior (Max=7 instead of Error=1) ---
+; RUN: not llvm-as %t/wrong-behavior.ll --disable-output 2>&1 \
+; RUN: | FileCheck %s --check-prefix=WRONG-BEHAVIOR
+
+; --- Negative: non-integer value ---
+; RUN: not llvm-as %t/non-integer.ll --disable-output 2>&1 \
+; RUN: | FileCheck %s --check-prefix=NON-INT
+
+; --- Negative: missing value ---
+; RUN: not llvm-as %t/missing-value.ll --disable-output 2>&1 \
+; RUN: | FileCheck %s --check-prefix=MISSING-VALUE
+
+; --- Negative: value out of range (2 is not 0 or 1) ---
+; RUN: not llvm-as %t/out-of-range.ll --disable-output 2>&1 \
+; RUN: | FileCheck %s --check-prefix=RANGE
+
+; WRONG-BEHAVIOR: 'amdgpu.xnack' module flag must use 'error' merge behaviour
+; NON-INT: 'amdgpu.xnack' module flag must have a constant integer value
+; MISSING-VALUE: incorrect number of operands in module flag
+; RANGE: 'amdgpu.xnack' module flag must be 0 or 1
+
+;--- wrong-behavior.ll
+; Max (i32 7) is not Error (i32 1).
+!0 = !{i32 7, !"amdgpu.xnack", i32 1}
+!llvm.module.flags = !{!0}
+
+;--- non-integer.ll
+; Error behavior but float value instead of integer.
+!0 = !{i32 1, !"amdgpu.xnack", float 1.0}
+!llvm.module.flags = !{!0}
+
+;--- missing-value.ll
+; Missing value field.
+!0 = !{i32 1, !"amdgpu.xnack"}
+!llvm.module.flags = !{!0}
+
+;--- out-of-range.ll
+; Value 2 is out of range (must be 0 or 1).
+!0 = !{i32 1, !"amdgpu.xnack", i32 2}
+!llvm.module.flags = !{!0}
More information about the llvm-branch-commits
mailing list