[clang] 27eeb73 - AMDGPU: Use module flags to control xnack and sramecc (#204595)
via cfe-commits
cfe-commits at lists.llvm.org
Tue Jul 21 05:38:44 PDT 2026
Author: Matt Arsenault
Date: 2026-07-21T14:38:39+02:00
New Revision: 27eeb7370281e4fdefdddeaa4b3cbb8ca11ea71f
URL: https://github.com/llvm/llvm-project/commit/27eeb7370281e4fdefdddeaa4b3cbb8ca11ea71f
DIFF: https://github.com/llvm/llvm-project/commit/27eeb7370281e4fdefdddeaa4b3cbb8ca11ea71f.diff
LOG: AMDGPU: Use module flags to control xnack and sramecc (#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>
Added:
clang/test/CodeGenOpenCL/amdgpu-module-flag-xnack-sramecc.cl
clang/test/CodeGenOpenCL/amdgpu-xnack-any-only.cl
llvm/test/CodeGen/AMDGPU/mattr-xnack-sramecc-legacy.ll
llvm/test/CodeGen/AMDGPU/module-flag-sramecc.ll
llvm/test/CodeGen/AMDGPU/module-flag-xnack-no-on-off-modes.ll
llvm/test/CodeGen/AMDGPU/module-flag-xnack-sramecc-combined.ll
llvm/test/CodeGen/AMDGPU/module-flag-xnack.ll
llvm/test/Linker/Inputs/amdgpu-sramecc-module-flag-0.ll
llvm/test/Linker/Inputs/amdgpu-sramecc-module-flag-1.ll
llvm/test/Linker/Inputs/amdgpu-sramecc-module-flag-any.ll
llvm/test/Linker/Inputs/amdgpu-xnack-module-flag-0.ll
llvm/test/Linker/Inputs/amdgpu-xnack-module-flag-1.ll
llvm/test/Linker/Inputs/amdgpu-xnack-module-flag-any.ll
llvm/test/Linker/amdgpu-sramecc-module-flag-0.ll
llvm/test/Linker/amdgpu-sramecc-module-flag-1.ll
llvm/test/Linker/amdgpu-sramecc-module-flag-any.ll
llvm/test/Linker/amdgpu-xnack-module-flag-0.ll
llvm/test/Linker/amdgpu-xnack-module-flag-1.ll
llvm/test/Linker/amdgpu-xnack-module-flag-any.ll
llvm/test/Verifier/AMDGPU/module-flag-sramecc.ll
llvm/test/Verifier/AMDGPU/module-flag-xnack.ll
Modified:
clang/include/clang/Basic/TargetOptions.h
clang/include/clang/CIR/Dialect/IR/CIRDialect.td
clang/include/clang/Options/Options.td
clang/lib/Basic/Targets/AMDGPU.cpp
clang/lib/CIR/CodeGen/CIRGenAMDGPU.cpp
clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp
clang/lib/CodeGen/CodeGenModule.cpp
clang/lib/Driver/ToolChains/Clang.cpp
clang/lib/Driver/ToolChains/Clang.h
clang/lib/Frontend/CompilerInvocation.cpp
clang/test/CIR/CodeGenHIP/target-features.hip
clang/test/Driver/amdgpu-features.c
clang/test/Driver/amdgpu-openmp-toolchain.c
clang/test/Driver/amdgpu-toolchain.c
clang/test/Driver/amdgpu-xnack-sramecc-flags.c
clang/test/Driver/hip-sanitize-options.hip
clang/test/Driver/hip-target-id.hip
clang/test/Driver/hip-toolchain-features.hip
clang/test/Driver/target-id.cl
llvm/docs/AMDGPUUsage.rst
llvm/docs/ReleaseNotes.md
llvm/lib/IR/VerifierAMDGPU.cpp
llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.h
llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
llvm/lib/Target/AMDGPU/GCNSubtarget.cpp
llvm/lib/Target/AMDGPU/GCNSubtarget.h
llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.cpp
llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.h
llvm/test/CodeGen/AMDGPU/GlobalISel/extractelement-stack-lower.ll
llvm/test/CodeGen/AMDGPU/amdpal-callable.ll
llvm/test/CodeGen/AMDGPU/break-smem-soft-clauses.mir
llvm/test/CodeGen/AMDGPU/callee-special-input-vgprs-packed.ll
llvm/test/CodeGen/AMDGPU/cluster-flat-loads-postra.mir
llvm/test/CodeGen/AMDGPU/cluster_stores.ll
llvm/test/CodeGen/AMDGPU/directive-amdgcn-target-legacy-triples.ll
llvm/test/CodeGen/AMDGPU/elf-header-flags-sramecc.ll
llvm/test/CodeGen/AMDGPU/elf-header-flags-xnack.ll
llvm/test/CodeGen/AMDGPU/flat-saddr-load.ll
llvm/test/CodeGen/AMDGPU/flat-scratch-reg.ll
llvm/test/CodeGen/AMDGPU/gfx902-without-xnack.ll
llvm/test/CodeGen/AMDGPU/greedy-reverse-local-assignment.ll
llvm/test/CodeGen/AMDGPU/hazard-hidden-bundle.mir
llvm/test/CodeGen/AMDGPU/hazard-in-bundle.mir
llvm/test/CodeGen/AMDGPU/hsa-metadata-kernel-code-props.ll
llvm/test/CodeGen/AMDGPU/hsa-metadata-resource-usage-function-ordering.ll
llvm/test/CodeGen/AMDGPU/hsa-note-no-func.ll
llvm/test/CodeGen/AMDGPU/immv216.ll
llvm/test/CodeGen/AMDGPU/limit-soft-clause-reg-pressure.mir
llvm/test/CodeGen/AMDGPU/materialize-frame-index-sgpr.ll
llvm/test/CodeGen/AMDGPU/nsa-reassign.ll
llvm/test/CodeGen/AMDGPU/nsa-vmem-hazard.mir
llvm/test/CodeGen/AMDGPU/occupancy-levels.ll
llvm/test/CodeGen/AMDGPU/post-ra-soft-clause-dbg-info.ll
llvm/test/CodeGen/AMDGPU/s_addk_i32.ll
llvm/test/CodeGen/AMDGPU/s_mulk_i32.ll
llvm/test/CodeGen/AMDGPU/schedule-amdgpu-tracker-physreg-crash.ll
llvm/test/CodeGen/AMDGPU/soft-clause-dbg-value.mir
llvm/test/CodeGen/AMDGPU/spill-scavenge-offset.ll
llvm/test/CodeGen/AMDGPU/sram-ecc-default.ll
llvm/test/CodeGen/AMDGPU/sramecc-subtarget-feature-disabled.ll
llvm/test/CodeGen/AMDGPU/sramecc-subtarget-feature-enabled.ll
llvm/test/CodeGen/AMDGPU/target-id-xnack-always-on.ll
llvm/test/CodeGen/AMDGPU/tid-kd-xnack-off.ll
llvm/test/CodeGen/AMDGPU/tid-kd-xnack-on.ll
llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-all-off.ll
llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-all-on.ll
llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-any-off-1.ll
llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-any-off-2.ll
llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-any-on-1.ll
llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-any-on-2.ll
llvm/test/CodeGen/AMDGPU/tid-one-func-xnack-off.ll
llvm/test/CodeGen/AMDGPU/tid-one-func-xnack-on.ll
llvm/test/CodeGen/AMDGPU/xnack-subtarget-feature-disabled.ll
llvm/test/CodeGen/AMDGPU/xnack-subtarget-feature-enabled.ll
llvm/test/MC/AMDGPU/xnack-mask.s
Removed:
llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-invalid-any-off-on.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/CIR/Dialect/IR/CIRDialect.td b/clang/include/clang/CIR/Dialect/IR/CIRDialect.td
index c20af04f97a1a..d340a9061310f 100644
--- a/clang/include/clang/CIR/Dialect/IR/CIRDialect.td
+++ b/clang/include/clang/CIR/Dialect/IR/CIRDialect.td
@@ -82,6 +82,8 @@ def CIR_Dialect : Dialect {
static llvm::StringRef getAMDGPUCodeObjectVersionAttrName() { return "cir.amdhsa_code_object_version"; }
static llvm::StringRef getAMDGPUPrintfKindAttrName() { return "cir.amdgpu_printf_kind"; }
+ static llvm::StringRef getAMDGPUXnackAttrName() { return "cir.amdgpu_xnack"; }
+ static llvm::StringRef getAMDGPUSramEccAttrName() { return "cir.amdgpu_sramecc"; }
static llvm::StringRef getOpenCLKernelArgMetadataAttrName() { return "cir.cl.kernel_arg_metadata"; }
void registerAttributes();
diff --git a/clang/include/clang/Options/Options.td b/clang/include/clang/Options/Options.td
index 07fec9e9aac21..c306179a084d3 100644
--- a/clang/include/clang/Options/Options.td
+++ b/clang/include/clang/Options/Options.td
@@ -5979,10 +5979,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 f32560c5b152b..c789fd8f94afb 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/CIR/CodeGen/CIRGenAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenAMDGPU.cpp
index 896e74e548c61..48a30edca3f54 100644
--- a/clang/lib/CIR/CodeGen/CIRGenAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenAMDGPU.cpp
@@ -38,4 +38,26 @@ void CIRGenModule::emitAMDGPUMetadata() {
theModule->setAttr(cir::CIRDialect::getAMDGPUPrintfKindAttrName(),
builder.getStringAttr(printfKind));
}
+
+ // Emit xnack module flag.
+ if (target.getTargetOpts().AMDGPUXnackState !=
+ TargetOptions::AMDGPUFeatureState::Any) {
+ theModule->setAttr(cir::CIRDialect::getAMDGPUXnackAttrName(),
+ builder.getI32IntegerAttr(
+ target.getTargetOpts().AMDGPUXnackState ==
+ TargetOptions::AMDGPUFeatureState::Enabled
+ ? 1
+ : 0));
+ }
+
+ // Emit sramecc module flag.
+ if (target.getTargetOpts().AMDGPUSramEccState !=
+ TargetOptions::AMDGPUFeatureState::Any) {
+ theModule->setAttr(cir::CIRDialect::getAMDGPUSramEccAttrName(),
+ builder.getI32IntegerAttr(
+ target.getTargetOpts().AMDGPUSramEccState ==
+ TargetOptions::AMDGPUFeatureState::Enabled
+ ? 1
+ : 0));
+ }
}
diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp
index 36b46d7a9c666..dd95f09ee77ae 100644
--- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp
+++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp
@@ -131,6 +131,22 @@ class CIRDialectLLVMIRTranslationInterface
}
}
+ if (attribute.getName() == "cir.amdgpu_xnack") {
+ if (auto intAttr =
+ mlir::dyn_cast<mlir::IntegerAttr>(attribute.getValue())) {
+ llvmModule->addModuleFlag(llvm::Module::Error, "amdgpu.xnack",
+ static_cast<uint32_t>(intAttr.getInt()));
+ }
+ }
+
+ if (attribute.getName() == "cir.amdgpu_sramecc") {
+ if (auto intAttr =
+ mlir::dyn_cast<mlir::IntegerAttr>(attribute.getValue())) {
+ llvmModule->addModuleFlag(llvm::Module::Error, "amdgpu.sramecc",
+ static_cast<uint32_t>(intAttr.getInt()));
+ }
+ }
+
return mlir::success();
}
};
diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp
index 5f5fc4401bb4e..8ce4063a3b780 100644
--- a/clang/lib/CodeGen/CodeGenModule.cpp
+++ b/clang/lib/CodeGen/CodeGenModule.cpp
@@ -1220,6 +1220,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 fb24d0b877ca6..10a1c0e8dec3f 100644
--- a/clang/lib/Driver/ToolChains/Clang.cpp
+++ b/clang/lib/Driver/ToolChains/Clang.cpp
@@ -1538,6 +1538,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 {
@@ -1565,6 +1574,10 @@ void Clang::RenderTargetOptions(const llvm::Triple &EffectiveTriple,
AddAArch64TargetArgs(Args, CmdArgs);
break;
+ case llvm::Triple::amdgpu:
+ 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 56362e758a78b..6562cb3a9b135 100644
--- a/clang/lib/Frontend/CompilerInvocation.cpp
+++ b/clang/lib/Frontend/CompilerInvocation.cpp
@@ -5022,6 +5022,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,
@@ -5053,6 +5065,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 (
diff ering 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/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 adbf809a6d8a1..c56cf4c97eac9 100644
--- a/clang/test/Driver/amdgpu-toolchain.c
+++ b/clang/test/Driver/amdgpu-toolchain.c
@@ -29,7 +29,7 @@
// 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
@@ -37,7 +37,7 @@
// RUN: %clang -### --target=amdgpu9.0a-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
diff erent 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 b7889b1634404..b9ff07cbe68d0 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 cbe9d15074457..b3b2a06005377 100644
--- a/llvm/docs/AMDGPUUsage.rst
+++ b/llvm/docs/AMDGPUUsage.rst
@@ -1012,10 +1012,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
diff erently
+ (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 in
diff erent 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 94208f8e36d19..5f47ab6dd925f 100644
--- a/llvm/docs/ReleaseNotes.md
+++ b/llvm/docs/ReleaseNotes.md
@@ -75,6 +75,9 @@ Makes programs 10x faster by doing Special New Thing.
### Changes to the AMDGPU Backend
+* Replaced `xnack` and `sramecc` target features with `amdgpu.xnack`
+ and `amdgpu.sramecc` module flags.
+
### Changes to the ARM Backend
### Changes to the AVR Backend
diff --git a/llvm/lib/IR/VerifierAMDGPU.cpp b/llvm/lib/IR/VerifierAMDGPU.cpp
index a668d3135644e..9f2cd159ad60f 100644
--- a/llvm/lib/IR/VerifierAMDGPU.cpp
+++ b/llvm/lib/IR/VerifierAMDGPU.cpp
@@ -40,18 +40,35 @@ using namespace llvm;
void llvm::verifyAMDGPUModuleFlag(VerifierSupport &VS, const MDString *ID,
Module::ModFlagBehavior MFB,
const MDNode *Op) {
- if (ID->getString() != "amdgpu.buffer.oob.mode" &&
- ID->getString() != "amdgpu.tbuffer.oob.mode")
+ StringRef FlagName = ID->getString();
+ if (!FlagName.consume_front("amdgpu."))
return;
- Check(MFB == Module::Max,
- "'" + ID->getString() + "' module flag must use 'max' 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");
- Check(Value->getZExtValue() <= 2,
- "'" + ID->getString() + "' module flag must be 0, 1, or 2");
+ if (FlagName == "buffer.oob.mode" || FlagName == "tbuffer.oob.mode") {
+ Check(MFB == Module::Max,
+ "'" + ID->getString() +
+ "' module flag must use 'max' 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");
+ Check(Value->getZExtValue() <= 2,
+ "'" + ID->getString() + "' module flag must be 0, 1, or 2");
+ return;
+ }
+
+ if (FlagName == "xnack" || FlagName == "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");
+ Check(Value->getZExtValue() <= 1,
+ "'" + ID->getString() + "' module flag must be 0 or 1");
+ return;
+ }
}
// Verify that when a function has !reqd_work_group_size metadata, it also has
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp b/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
index a01130af415db..fec2b9a97dd7d 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
@@ -205,30 +205,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() != AMDGPU::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() != AMDGPU::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;
@@ -1205,31 +1181,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 AMDGPU::TargetID &STMTargetID = STM.getTargetID();
- if (TSTargetID->isXnackSupported())
- if (TSTargetID->getXnackSetting() == AMDGPU::TargetIDSetting::Any)
- TSTargetID->setXnackSetting(STMTargetID.getXnackSetting());
- if (TSTargetID->isSramEccSupported())
- if (TSTargetID->getSramEccSetting() == AMDGPU::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::TargetIDSetting Setting =
+ GCNTargetMachine::getTargetIDSettingFromModuleFlag(M, "amdgpu.xnack");
+ if (Setting != AMDGPU::TargetIDSetting::Any)
+ TSTargetID->setXnackSetting(Setting);
+ }
+
+ if (getGlobalSTI()->getFeatureBits().test(AMDGPU::FeatureSupportsSRAMECC)) {
+ AMDGPU::TargetIDSetting Setting =
+ GCNTargetMachine::getTargetIDSettingFromModuleFlag(M, "amdgpu.sramecc");
+ if (Setting != AMDGPU::TargetIDSetting::Any)
+ TSTargetID->setSramEccSetting(Setting);
}
}
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp b/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
index 122d3285e1d11..d1fb7a2ea1193 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
@@ -470,6 +470,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",
@@ -1300,6 +1310,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::TargetIDSetting
+GCNTargetMachine::getTargetIDSettingFromModuleFlag(const Module &M,
+ StringRef FlagName) {
+ using AMDGPU::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);
@@ -1310,12 +1340,26 @@ GCNTargetMachine::getSubtargetImpl(const Function &F) const {
OOBFlagValue TBufOOB = getOOBFlagValue(M, AMDGPUOOBMode::TBufferFlag);
bool BufRelaxed = BufOOB == OOBFlagValue::Relaxed;
bool TBufRelaxed = TBufOOB == OOBFlagValue::Relaxed;
+
+ using AMDGPU::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) {
@@ -1339,7 +1383,7 @@ GCNTargetMachine::getSubtargetImpl(const Function &F) const {
}
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..189b6399dbf7b 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::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 7ead3d6f2b263..e01307cd4b12b 100644
--- a/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
+++ b/llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
@@ -9569,7 +9569,7 @@ void AMDGPUAsmParser::onBeginOfFile() {
if (!getTargetStreamer().getTargetID())
getTargetStreamer().initializeTargetID(getSTI(),
- getSTI().getFeatureString());
+ /*ApplyFeatureString=*/true);
}
void AMDGPUAsmParser::emitTargetDirective() {
diff --git a/llvm/lib/Target/AMDGPU/GCNSubtarget.cpp b/llvm/lib/Target/AMDGPU/GCNSubtarget.cpp
index 03ead5b978d6c..1c0e718bd8d97 100644
--- a/llvm/lib/Target/AMDGPU/GCNSubtarget.cpp
+++ b/llvm/lib/Target/AMDGPU/GCNSubtarget.cpp
@@ -196,11 +196,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;
}
@@ -217,11 +212,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::TargetIDSetting XnackSetting,
+ AMDGPU::TargetIDSetting SramEccSetting)
: // clang-format off
AMDGPUGenSubtargetInfo(TT, GPU, /*TuneCPU*/ GPU, FS),
AMDGPUSubtarget(TT),
- TargetID(AMDGPU::createAMDGPUTargetID(*this, FS)),
+ TargetID(AMDGPU::createAMDGPUTargetID(*this, "")),
InstrItins(getInstrItineraryForCPU(GPU)),
BufferOOBRelaxed(BufferOOBRelaxed),
TBufferOOBRelaxed(TBufferOOBRelaxed),
@@ -232,6 +229,22 @@ GCNSubtarget::GCNSubtarget(const Triple &TT, StringRef GPU, StringRef FS,
/*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 170ec7f85aa23..ef00be6f1ab4e 100644
--- a/llvm/lib/Target/AMDGPU/GCNSubtarget.h
+++ b/llvm/lib/Target/AMDGPU/GCNSubtarget.h
@@ -111,9 +111,11 @@ class GCNSubtarget final : public AMDGPUGenSubtargetInfo,
const MachineInstr &UseI, int UseOpIdx) const;
public:
- GCNSubtarget(const Triple &TT, StringRef GPU, StringRef FS,
- const GCNTargetMachine &TM, bool BufferOOBRelaxed = false,
- bool TBufferOOBRelaxed = false);
+ GCNSubtarget(
+ const Triple &TT, StringRef GPU, StringRef FS, const GCNTargetMachine &TM,
+ bool BufferOOBRelaxed = false, bool TBufferOOBRelaxed = false,
+ AMDGPU::TargetIDSetting XnackSetting = AMDGPU::TargetIDSetting::Any,
+ AMDGPU::TargetIDSetting SramEccSetting = AMDGPU::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 d70ad37cca867..8f3789ea4d65f 100644
--- a/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.cpp
+++ b/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.cpp
@@ -46,6 +46,15 @@ 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 = AMDGPU::createAMDGPUTargetID(
+ 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 c42ec5c2683cf..d1ef710e9b180 100644
--- a/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.h
+++ b/llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.h
@@ -137,10 +137,8 @@ class AMDGPUTargetStreamer : public MCTargetStreamer {
return TargetID;
}
std::optional<AMDGPU::TargetID> &getTargetID() { return TargetID; }
- void initializeTargetID(const MCSubtargetInfo &STI, StringRef FeatureString) {
- assert(TargetID == std::nullopt && "TargetID can only be initialized once");
- TargetID = AMDGPU::createAMDGPUTargetID(STI, FeatureString);
- }
+ void initializeTargetID(const MCSubtargetInfo &STI,
+ bool ApplyFeatureString = false);
};
class AMDGPUTargetAsmStreamer final : public AMDGPUTargetStreamer {
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/extractelement-stack-lower.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/extractelement-stack-lower.ll
index 6621453083c1d..ce280584cbbe3 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 -mtriple=amdgpu9.00-mesa-mesa3d -mattr=-xnack < %s | FileCheck -check-prefixes=GFX9 %s
+; RUN: llc -global-isel -mtriple=amdgpu9.00-mesa-mesa3d --amdgpu-xnack=false < %s | FileCheck -check-prefixes=GFX9 %s
; RUN: llc -global-isel -mtriple=amdgpu12.00-mesa-mesa3d < %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 0ca1680ed36f4..20e9842d5e62b 100644
--- a/llvm/test/CodeGen/AMDGPU/amdpal-callable.ll
+++ b/llvm/test/CodeGen/AMDGPU/amdpal-callable.ll
@@ -1,6 +1,6 @@
-; RUN: llc -mtriple=amdgpu6.00--amdpal -mattr=-xnack -mattr=+dx10-clamp-and-ieee-mode < %s | FileCheck -check-prefixes=GCN,SDAG,GFX8 -enable-var-scope %s
-; RUN: llc -mtriple=amdgpu9.00--amdpal -mattr=-xnack < %s | FileCheck -check-prefixes=GCN,SDAG,GFX9 -enable-var-scope %s
-; RUN: llc -global-isel -mtriple=amdgpu9.00--amdpal -mattr=-xnack < %s | FileCheck -check-prefixes=GCN,GISEL,GFX9 -enable-var-scope %s
+; RUN: llc -mtriple=amdgpu6.00--amdpal --amdgpu-xnack=false -mattr=+dx10-clamp-and-ieee-mode < %s | FileCheck -check-prefixes=GCN,SDAG,GFX8 -enable-var-scope %s
+; RUN: llc -mtriple=amdgpu9.00--amdpal --amdgpu-xnack=false < %s | FileCheck -check-prefixes=GCN,SDAG,GFX9 -enable-var-scope %s
+; RUN: llc -global-isel -mtriple=amdgpu9.00--amdpal --amdgpu-xnack=false < %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 f5116a808d950..8c5ef9cd5525b 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=amdgpu8.01 -verify-machineinstrs -run-pass post-RA-hazard-rec %s -o - | FileCheck -check-prefixes=GCN,XNACK %s
-# RUN: llc -mtriple=amdgpu8.03 -mattr=-xnack -verify-machineinstrs -run-pass post-RA-hazard-rec %s -o - | FileCheck -check-prefixes=GCN %s
+# RUN: llc -mtriple=amdgpu8.03 --amdgpu-xnack=false -verify-machineinstrs -run-pass post-RA-hazard-rec %s -o - | FileCheck -check-prefixes=GCN %s
-# RUN: llc -mtriple=amdgpu8.03 -mattr=-xnack -passes post-RA-hazard-rec %s -o - | FileCheck -check-prefixes=GCN %s
+# RUN: llc -mtriple=amdgpu8.03 --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 6396ea498853a..03525a359de91 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 -mtriple=amdgpu7.00-amd-amdhsa -passes=amdgpu-attributor < %s | llc -enable-ipra=0 | FileCheck -enable-var-scope -check-prefixes=GCN,GFX7 %s
-; RUN: opt -mtriple=amdgpu9.0a-amd-amdhsa -passes=amdgpu-attributor -mattr=-xnack < %s | llc -mattr=-xnack -enable-ipra=0 | FileCheck -enable-var-scope -check-prefixes=GCN,GFX90A %s
+; RUN: opt -mtriple=amdgpu9.0a-amd-amdhsa -passes=amdgpu-attributor --amdgpu-xnack=false < %s | llc --amdgpu-xnack=false -enable-ipra=0 | FileCheck -enable-var-scope -check-prefixes=GCN,GFX90A %s
define void @use_workitem_id_x() #1 {
; GFX7-LABEL: use_workitem_id_x:
diff --git a/llvm/test/CodeGen/AMDGPU/cluster-flat-loads-postra.mir b/llvm/test/CodeGen/AMDGPU/cluster-flat-loads-postra.mir
index c2d50a83a5888..9e4e2aea19203 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=amdgpu8.02 -mattr=-xnack -run-pass post-RA-sched -verify-machineinstrs -o - %s | FileCheck -check-prefix=GCN %s
-# RUN: llc -mtriple=amdgpu8.02 -mattr=-xnack -passes=post-RA-sched -o - %s | FileCheck -check-prefix=GCN %s
+# RUN: llc -mtriple=amdgpu8.02 --amdgpu-xnack=false -run-pass post-RA-sched -verify-machineinstrs -o - %s | FileCheck -check-prefix=GCN %s
+# RUN: llc -mtriple=amdgpu8.02 --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 43340d42e1628..0aaa8167b5fe2 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=amdgpu9.00 -mattr=-xnack -debug-only=machine-scheduler < %s 2> %t | FileCheck --enable-var-scope --check-prefix=GFX9 %s
+; RUN: llc -mtriple=amdgpu9.00 --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=amdgpu10.10 -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-legacy-triples.ll b/llvm/test/CodeGen/AMDGPU/directive-amdgcn-target-legacy-triples.ll
index 0164332637131..92b52f025ea13 100644
--- a/llvm/test/CodeGen/AMDGPU/directive-amdgcn-target-legacy-triples.ll
+++ b/llvm/test/CodeGen/AMDGPU/directive-amdgcn-target-legacy-triples.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
@@ -114,12 +114,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=gfx11-7-generic < %s | FileCheck --check-prefixes=GFX11_7_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 814fbc8aaa69b..92d729aab9295 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=amdgpu9.06 < %s | llvm-readobj --file-headers - | FileCheck --check-prefixes=SRAM-ECC-GFX906 %s
-; RUN: llc -filetype=obj -mtriple=amdgpu9.06 -mattr=-sramecc < %s | llvm-readobj --file-headers - | FileCheck --check-prefixes=NO-SRAM-ECC-GFX906 %s
-; RUN: llc -filetype=obj -mtriple=amdgpu9.06 -mattr=+sramecc < %s | llvm-readobj --file-headers - | FileCheck --check-prefixes=SRAM-ECC-GFX906 %s
-; RUN: llc -filetype=obj -mtriple=amdgpu9.06 -mattr=+sramecc,+xnack < %s | llvm-readobj --file-headers - | FileCheck --check-prefixes=SRAM-ECC-XNACK-GFX906 %s
+; RUN: llc -filetype=obj -mtriple=amdgpu9.06 -amdgpu-sramecc=0 < %s | llvm-readobj --file-headers - | FileCheck --check-prefixes=NO-SRAM-ECC-GFX906 %s
+; RUN: llc -filetype=obj -mtriple=amdgpu9.06 -amdgpu-sramecc=1 < %s | llvm-readobj --file-headers - | FileCheck --check-prefixes=SRAM-ECC-GFX906 %s
+; RUN: llc -filetype=obj -mtriple=amdgpu9.06 -amdgpu-sramecc=1 -amdgpu-xnack=1 < %s | llvm-readobj --file-headers - | FileCheck --check-prefixes=SRAM-ECC-XNACK-GFX906 %s
; RUN: llc -filetype=obj -mtriple=amdgpu9.08 < %s | llvm-readobj --file-header - | FileCheck --check-prefix=SRAM-ECC-GFX908 %s
-; RUN: llc -filetype=obj -mtriple=amdgpu9.08 -mattr=+sramecc < %s | llvm-readobj --file-header - | FileCheck --check-prefix=SRAM-ECC-GFX908 %s
+; RUN: llc -filetype=obj -mtriple=amdgpu9.08 -amdgpu-sramecc=1 < %s | llvm-readobj --file-header - | FileCheck --check-prefix=SRAM-ECC-GFX908 %s
; RUN: llc -filetype=obj -mtriple=amdgpu9.0a < %s | llvm-readobj --file-header - | FileCheck --check-prefix=SRAM-ECC-GFX90A %s
-; RUN: llc -filetype=obj -mtriple=amdgpu9.0a -mattr=+sramecc < %s | llvm-readobj --file-header - | FileCheck --check-prefix=SRAM-ECC-GFX90A %s
+; RUN: llc -filetype=obj -mtriple=amdgpu9.0a -amdgpu-sramecc=1 < %s | llvm-readobj --file-header - | FileCheck --check-prefix=SRAM-ECC-GFX90A %s
; RUN: llc -filetype=obj -mtriple=amdgpu9.42 < %s | llvm-readobj --file-header - | FileCheck --check-prefix=SRAM-ECC-GFX942 %s
-; RUN: llc -filetype=obj -mtriple=amdgpu9.42 -mattr=+sramecc < %s | llvm-readobj --file-header - | FileCheck --check-prefix=SRAM-ECC-GFX942 %s
+; RUN: llc -filetype=obj -mtriple=amdgpu9.42 -amdgpu-sramecc=1 < %s | llvm-readobj --file-header - | FileCheck --check-prefix=SRAM-ECC-GFX942 %s
; RUN: llc -filetype=obj -mtriple=amdgpu9.50 < %s | llvm-readobj --file-header - | FileCheck --check-prefix=SRAM-ECC-GFX950 %s
-; RUN: llc -filetype=obj -mtriple=amdgpu9.50 -mattr=+sramecc < %s | llvm-readobj --file-header - | FileCheck --check-prefix=SRAM-ECC-GFX950 %s
+; RUN: llc -filetype=obj -mtriple=amdgpu9.50 -amdgpu-sramecc=1 < %s | llvm-readobj --file-header - | FileCheck --check-prefix=SRAM-ECC-GFX950 %s
; RUN: llc -filetype=obj -mtriple=amdgpu12.50 < %s | llvm-readobj --file-header - | FileCheck --check-prefix=SRAM-ECC-GFX1250 %s
-; RUN: llc -filetype=obj -mtriple=amdgpu12.50 -mattr=+sramecc < %s | llvm-readobj --file-header - | FileCheck --check-prefix=SRAM-ECC-GFX1250 %s
+; RUN: llc -filetype=obj -mtriple=amdgpu12.50 -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 97f85cfe13c67..89bc7cf8f6546 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=amdgpu8.01 < %s | llvm-readobj --file-headers - | FileCheck --check-prefixes=XNACK-GFX801 %s
-; RUN: llc -filetype=obj -mtriple=amdgpu8.01 -mattr=+xnack < %s | llvm-readobj --file-headers - | FileCheck --check-prefixes=XNACK-GFX801 %s
+; RUN: llc -filetype=obj -mtriple=amdgpu8.01 --amdgpu-xnack=true < %s | llvm-readobj --file-headers - | FileCheck --check-prefixes=XNACK-GFX801 %s
; RUN: llc -filetype=obj -mtriple=amdgpu8.02 < %s | llvm-readobj --file-headers - | FileCheck --check-prefixes=NO-XNACK-GFX802 %s
-; RUN: llc -filetype=obj -mtriple=amdgpu8.02 -mattr=-xnack < %s | llvm-readobj --file-headers - | FileCheck --check-prefixes=NO-XNACK-GFX802 %s
+; RUN: llc -filetype=obj -mtriple=amdgpu8.02 --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 df17904c3556d..f1b6f0879fce5 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=amdgpu12.50-mesa-mesa3d -mattr=+real-true16 < %s | FileCheck -check-prefixes=GFX1250,GFX1250-SDAG,GFX1250-SDAG-TRUE16 %s
; RUN: llc -global-isel=1 -mtriple=amdgpu12.50-mesa-mesa3d -mattr=+real-true16 < %s | FileCheck -check-prefixes=GFX1250,GFX1250-GISEL,GFX1250-GISEL-TRUE16 %s
-; RUN: llc -mtriple=amdgpu12.50-mesa-mesa3d -mattr=+real-true16,-sramecc < %s | FileCheck -check-prefixes=GFX1250,GFX1250-NOECC,GFX1250-NOECC-SDAG-TRUE16 %s
-; RUN: llc -mtriple=amdgpu12.50-mesa-mesa3d -mattr=-real-true16,-sramecc < %s | FileCheck -check-prefixes=GFX1250,GFX1250-NOECC,GFX1250-NOECC-SDAG-FAKE16 %s
+; RUN: llc -mtriple=amdgpu12.50-mesa-mesa3d -mattr=+real-true16 --amdgpu-sramecc=false < %s | FileCheck -check-prefixes=GFX1250,GFX1250-NOECC,GFX1250-NOECC-SDAG-TRUE16 %s
+; RUN: llc -mtriple=amdgpu12.50-mesa-mesa3d -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 9de6c371d8766..b4f39cfa7e96b 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=amdgpu7.00 | FileCheck -check-prefix=CI -check-prefix=GCN %s
-; RUN: llc < %s -mtriple=amdgpu8.03 -mattr=-xnack | FileCheck -check-prefix=VI-NOXNACK -check-prefix=GCN %s
+; RUN: llc < %s -mtriple=amdgpu8.03 --amdgpu-xnack=false | FileCheck -check-prefix=VI-NOXNACK -check-prefix=GCN %s
-; RUN: llc < %s -mtriple=amdgpu8.01 -mattr=-xnack | FileCheck -check-prefixes=VI-NOXNACK,GCN %s
-; RUN: llc < %s -mtriple=amdgpu8.10 -mattr=-xnack | FileCheck -check-prefixes=VI-NOXNACK,GCN %s
+; RUN: llc < %s -mtriple=amdgpu8.01 --amdgpu-xnack=false | FileCheck -check-prefixes=VI-NOXNACK,GCN %s
+; RUN: llc < %s -mtriple=amdgpu8.10 --amdgpu-xnack=false | FileCheck -check-prefixes=VI-NOXNACK,GCN %s
-; RUN: llc < %s -mtriple=amdgpu8.01 -mattr=+xnack | FileCheck -check-prefix=VI-XNACK -check-prefix=GCN %s
-; RUN: llc < %s -mtriple=amdgpu8.10 -mattr=+xnack | FileCheck -check-prefix=VI-XNACK -check-prefix=GCN %s
+; RUN: llc < %s -mtriple=amdgpu8.01 --amdgpu-xnack=true | FileCheck -check-prefix=VI-XNACK -check-prefix=GCN %s
+; RUN: llc < %s -mtriple=amdgpu8.10 --amdgpu-xnack=true | FileCheck -check-prefix=VI-XNACK -check-prefix=GCN %s
; RUN: llc < %s -mtriple=amdgpu7.00--amdhsa | FileCheck -check-prefixes=GCN %s
-; RUN: llc < %s -mtriple=amdgpu8.01--amdhsa -mattr=-xnack | FileCheck -check-prefixes=HSA-VI-NOXNACK,GCN %s
-; RUN: llc < %s -mtriple=amdgpu8.01--amdhsa -mattr=+xnack | FileCheck -check-prefixes=HSA-VI-XNACK,GCN %s
+; RUN: llc < %s -mtriple=amdgpu8.01--amdhsa --amdgpu-xnack=false | FileCheck -check-prefixes=HSA-VI-NOXNACK,GCN %s
+; RUN: llc < %s -mtriple=amdgpu8.01--amdhsa --amdgpu-xnack=true | FileCheck -check-prefixes=HSA-VI-XNACK,GCN %s
; RUN: llc < %s -mtriple=amdgpu9.00--amdhsa -mattr=+architected-flat-scratch | FileCheck -check-prefixes=GCN %s
-; RUN: llc < %s -mtriple=amdgpu9.00--amdhsa -mattr=+architected-flat-scratch,-xnack | FileCheck -check-prefixes=GFX9-ARCH-FLAT-NOXNACK,GCN %s
-; RUN: llc < %s -mtriple=amdgpu9.00--amdhsa -mattr=+architected-flat-scratch,+xnack | FileCheck -check-prefixes=GFX9-ARCH-FLAT-XNACK,GCN %s
+; RUN: llc < %s -mtriple=amdgpu9.00--amdhsa -mattr=+architected-flat-scratch --amdgpu-xnack=false | FileCheck -check-prefixes=GFX9-ARCH-FLAT-NOXNACK,GCN %s
+; RUN: llc < %s -mtriple=amdgpu9.00--amdhsa -mattr=+architected-flat-scratch --amdgpu-xnack=true | FileCheck -check-prefixes=GFX9-ARCH-FLAT-XNACK,GCN %s
; RUN: llc < %s -mtriple=amdgpu10.10--amdhsa -mattr=+architected-flat-scratch | FileCheck -check-prefixes=GCN %s
-; RUN: llc < %s -mtriple=amdgpu10.10--amdhsa -mattr=+architected-flat-scratch,-xnack | FileCheck -check-prefixes=GFX10-ARCH-FLAT-NOXNACK,GCN %s
-; RUN: llc < %s -mtriple=amdgpu10.10--amdhsa -mattr=+architected-flat-scratch,+xnack | FileCheck -check-prefixes=GFX10-ARCH-FLAT-XNACK,GCN %s
+; RUN: llc < %s -mtriple=amdgpu10.10--amdhsa -mattr=+architected-flat-scratch --amdgpu-xnack=false | FileCheck -check-prefixes=GFX10-ARCH-FLAT-NOXNACK,GCN %s
+; RUN: llc < %s -mtriple=amdgpu10.10--amdhsa -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 cf6ab87e32ab7..706a236dfbf40 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=amdgpu9.02-amd-amdhsa -mattr=-xnack < %s | FileCheck %s
+; RUN: llc -mtriple=amdgpu9.02-amd-amdhsa < %s | FileCheck %s
; CHECK: .amdgcn_target "amdgpu9.02-amd-amdhsa-unknown-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 8440ad85f4be9..2ec0ca5afc586 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=amdgpu9.00-amd-amdhsa -greedy-reverse-local-assignment=0 < %s | FileCheck -check-prefixes=FORWARDXNACK %s
; RUN: llc -mtriple=amdgpu9.00-amd-amdhsa -greedy-reverse-local-assignment=1 < %s | FileCheck -check-prefixes=REVERSEXNACK %s
-; RUN: llc -mtriple=amdgpu9.00-amd-amdhsa -greedy-reverse-local-assignment=0 -mattr=-xnack < %s | FileCheck -check-prefix=NOXNACK %s
-; RUN: llc -mtriple=amdgpu9.00-amd-amdhsa -greedy-reverse-local-assignment=1 -mattr=-xnack < %s | FileCheck -check-prefix=NOXNACK %s
+; RUN: llc -mtriple=amdgpu9.00-amd-amdhsa -greedy-reverse-local-assignment=0 --amdgpu-xnack=false < %s | FileCheck -check-prefix=NOXNACK %s
+; RUN: llc -mtriple=amdgpu9.00-amd-amdhsa -greedy-reverse-local-assignment=1 --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 e099a1a5b485c..8908d2b982b85 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=amdgpu9.02 -mattr=+xnack -verify-machineinstrs -run-pass post-RA-hazard-rec %s -o - | FileCheck -check-prefixes=GCN,XNACK,GFX9 %s
-# RUN: llc -mtriple=amdgpu9.00 -mattr=-xnack -verify-machineinstrs -run-pass post-RA-hazard-rec %s -o - | FileCheck -check-prefixes=GCN,NOXNACK,GFX9 %s
-# RUN: llc -mtriple=amdgpu10.10 -mattr=+wavefrontsize64,-xnack -verify-machineinstrs -run-pass post-RA-hazard-rec %s -o - | FileCheck -check-prefixes=GCN,NOXNACK %s
+# RUN: llc -mtriple=amdgpu9.02 --amdgpu-xnack=true -verify-machineinstrs -run-pass post-RA-hazard-rec %s -o - | FileCheck -check-prefixes=GCN,XNACK,GFX9 %s
+# RUN: llc -mtriple=amdgpu9.00 --amdgpu-xnack=false -verify-machineinstrs -run-pass post-RA-hazard-rec %s -o - | FileCheck -check-prefixes=GCN,NOXNACK,GFX9 %s
+# RUN: llc -mtriple=amdgpu10.10 -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 f7e358f25d316..3cc5d455d36d8 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=amdgpu9.02 -mattr=+xnack -verify-machineinstrs -run-pass post-RA-hazard-rec %s -o - | FileCheck -check-prefixes=GCN,XNACK,GFX9 %s
-# RUN: llc -mtriple=amdgpu9.00 -mattr=-xnack -verify-machineinstrs -run-pass post-RA-hazard-rec %s -o - | FileCheck -check-prefixes=GCN,NOXNACK,GFX9 %s
-# RUN: llc -mtriple=amdgpu10.10 -mattr=+wavefrontsize64,-xnack -verify-machineinstrs -run-pass post-RA-hazard-rec %s -o - | FileCheck -check-prefixes=GCN,NOXNACK,GFX10 %s
+# RUN: llc -mtriple=amdgpu9.02 --amdgpu-xnack=true -verify-machineinstrs -run-pass post-RA-hazard-rec %s -o - | FileCheck -check-prefixes=GCN,XNACK,GFX9 %s
+# RUN: llc -mtriple=amdgpu9.00 --amdgpu-xnack=false -verify-machineinstrs -run-pass post-RA-hazard-rec %s -o - | FileCheck -check-prefixes=GCN,NOXNACK,GFX9 %s
+# RUN: llc -mtriple=amdgpu10.10 -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 9e092dcd8f559..06ecf2cc49402 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=amdgpu7.00-amd-amdhsa -enable-misched=0 -filetype=obj -o - < %s | llvm-readelf --notes - | FileCheck --check-prefixes=CHECK,GFX700,WAVE64 %s
-; RUN: llc -mattr=-xnack -mtriple=amdgpu8.03-amd-amdhsa -enable-misched=0 -filetype=obj -o - < %s | llvm-readelf --notes - | FileCheck --check-prefixes=CHECK,GFX803,WAVE64 %s
-; RUN: llc -mattr=-xnack -mtriple=amdgpu9.00-amd-amdhsa -enable-misched=0 -filetype=obj -o - < %s | llvm-readelf --notes - | FileCheck --check-prefixes=CHECK,GFX900,WAVE64 %s
-; RUN: llc -mattr=-xnack -mtriple=amdgpu10.10-amd-amdhsa -enable-misched=0 -filetype=obj -o - < %s | llvm-readelf --notes - | FileCheck --check-prefixes=CHECK,GFX1010,WAVE32 %s
+; RUN: llc --amdgpu-xnack=false -mtriple=amdgpu8.03-amd-amdhsa -enable-misched=0 -filetype=obj -o - < %s | llvm-readelf --notes - | FileCheck --check-prefixes=CHECK,GFX803,WAVE64 %s
+; RUN: llc --amdgpu-xnack=false -mtriple=amdgpu9.00-amd-amdhsa -enable-misched=0 -filetype=obj -o - < %s | llvm-readelf --notes - | FileCheck --check-prefixes=CHECK,GFX900,WAVE64 %s
+; RUN: llc --amdgpu-xnack=false -mtriple=amdgpu10.10-amd-amdhsa -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 827d287ab70eb..9b7c84f156066 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=amdgpu7.00-amd-amdhsa -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=amdgpu8.03-amd-amdhsa -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=amdgpu9.00-amd-amdhsa -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=amdgpu10.10-amd-amdhsa -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=amdgpu8.03-amd-amdhsa -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=amdgpu9.00-amd-amdhsa -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=amdgpu10.10-amd-amdhsa -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 59816af7aa6ba..99c3e8f966916 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=amdgpu8.05--amdhsa | FileCheck --check-prefix=HSA-VI805 %s
; RUN: llc < %s -mtriple=amdgpu8.10--amdhsa | FileCheck --check-prefix=HSA-VI810 %s
; RUN: llc < %s -mtriple=amdgpu8.10--amdhsa | FileCheck --check-prefix=HSA-VI810 %s
-; RUN: llc < %s -mtriple=amdgpu9.00--amdhsa -mattr=-xnack | FileCheck --check-prefix=HSA-GFX900 %s
+; RUN: llc < %s -mtriple=amdgpu9.00--amdhsa --amdgpu-xnack=false | FileCheck --check-prefix=HSA-GFX900 %s
; RUN: llc < %s -mtriple=amdgpu9.00--amdhsa | FileCheck --check-prefix=HSA-GFX901 %s
-; RUN: llc < %s -mtriple=amdgpu9.02--amdhsa -mattr=-xnack | FileCheck --check-prefix=HSA-GFX902 %s
+; RUN: llc < %s -mtriple=amdgpu9.02--amdhsa --amdgpu-xnack=false | FileCheck --check-prefix=HSA-GFX902 %s
; RUN: llc < %s -mtriple=amdgpu9.02--amdhsa | FileCheck --check-prefix=HSA-GFX903 %s
-; RUN: llc < %s -mtriple=amdgpu9.04--amdhsa -mattr=-xnack | FileCheck --check-prefix=HSA-GFX904 %s
+; RUN: llc < %s -mtriple=amdgpu9.04--amdhsa --amdgpu-xnack=false | FileCheck --check-prefix=HSA-GFX904 %s
; RUN: llc < %s -mtriple=amdgpu9.04--amdhsa | FileCheck --check-prefix=HSA-GFX905 %s
-; RUN: llc < %s -mtriple=amdgpu9.06--amdhsa -mattr=-xnack | FileCheck --check-prefix=HSA-GFX906 %s
+; RUN: llc < %s -mtriple=amdgpu9.06--amdhsa --amdgpu-xnack=false | FileCheck --check-prefix=HSA-GFX906 %s
; RUN: llc < %s -mtriple=amdgpu9.06--amdhsa | FileCheck --check-prefix=HSA-GFX907 %s
; NONHSA-SI600: .amd_amdgpu_isa "amdgpu6.00-unknown-unknown-unknown-gfx600"
diff --git a/llvm/test/CodeGen/AMDGPU/immv216.ll b/llvm/test/CodeGen/AMDGPU/immv216.ll
index dac9e2abcdd86..e6ae66b136ba1 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=amdgpu11.00--amdhsa -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=amdgpu10.10--amdhsa -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=amdgpu9.00--amdhsa -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=amdgpu8.03--amdhsa -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=amdgpu10.10--amdhsa -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=amdgpu9.00--amdhsa -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=amdgpu8.03--amdhsa -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=amdgpu7.00--amdhsa -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 51fe6eaba8450..57270ed42786d 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=amdgpu9.06-amd-amdhsa -mattr=+xnack -run-pass=si-form-memory-clauses -verify-machineinstrs -o - %s | FileCheck %s
-# RUN: llc -mtriple=amdgpu9.06-amd-amdhsa -mattr=+xnack -passes="si-form-memory-clauses" -o - %s | FileCheck %s
+# RUN: llc -mtriple=amdgpu9.06-amd-amdhsa --amdgpu-xnack=true -run-pass=si-form-memory-clauses -verify-machineinstrs -o - %s | FileCheck %s
+# RUN: llc -mtriple=amdgpu9.06-amd-amdhsa --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 42f5b113d967d..67065a6db8028 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=amdgpu7.00-amd-amdhsa < %s | FileCheck -check-prefix=GFX7 %s
-; RUN: llc -mtriple=amdgpu8.10-amd-amdhsa -mattr=+xnack < %s | FileCheck -check-prefix=GFX8 %s
-; RUN: llc -mtriple=amdgpu9.00-amd-amdhsa -mattr=+xnack < %s | FileCheck -check-prefixes=GFX900 %s
-; RUN: llc -mtriple=amdgpu9.42-amd-amdhsa -mattr=+xnack < %s | FileCheck -check-prefixes=GFX942 %s
+; RUN: llc -mtriple=amdgpu8.10-amd-amdhsa --amdgpu-xnack=true < %s | FileCheck -check-prefix=GFX8 %s
+; RUN: llc -mtriple=amdgpu9.00-amd-amdhsa --amdgpu-xnack=true < %s | FileCheck -check-prefixes=GFX900 %s
+; RUN: llc -mtriple=amdgpu9.42-amd-amdhsa --amdgpu-xnack=true < %s | FileCheck -check-prefixes=GFX942 %s
; RUN: llc -mtriple=amdgpu10.10-amd-amdhsa < %s | FileCheck -check-prefix=GFX10_1 %s
; RUN: llc -mtriple=amdgpu10.30-amd-amdhsa < %s | FileCheck -check-prefix=GFX10_3 %s
; RUN: llc -mtriple=amdgpu11.00-amd-amdhsa < %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..d58bb0567a35b
--- /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=amdgpu9.0a-amd-amdhsa -mattr=+xnack,+sramecc < %s 2>&1 | FileCheck --check-prefix=BOTH-ERR %s
+; RUN: not llc -mtriple=amdgpu9.0a-amd-amdhsa -mattr=+xnack < %s 2>&1 | FileCheck --check-prefix=XNACK-ERR %s
+; RUN: not llc -mtriple=amdgpu9.0a-amd-amdhsa -mattr=-xnack < %s 2>&1 | FileCheck --check-prefix=XNACK-ERR %s
+; RUN: not llc -mtriple=amdgpu9.0a-amd-amdhsa -mattr=xnack < %s 2>&1 | FileCheck --check-prefix=XNACK-ERR %s
+; RUN: not llc -mtriple=amdgpu9.0a-amd-amdhsa -mattr=+sramecc < %s 2>&1 | FileCheck --check-prefix=SRAMECC-ERR %s
+; RUN: not llc -mtriple=amdgpu9.0a-amd-amdhsa -mattr=-sramecc < %s 2>&1 | FileCheck --check-prefix=SRAMECC-ERR %s
+; RUN: not llc -mtriple=amdgpu9.0a-amd-amdhsa -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..ecb3266160441
--- /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=amdgpu9.06-amd-amdhsa < %t/on.ll | FileCheck --check-prefix=SRAMECC-ON %s
+; RUN: llc -mtriple=amdgpu9.06-amd-amdhsa < %t/off.ll | FileCheck --check-prefix=SRAMECC-OFF %s
+; RUN: llc -mtriple=amdgpu9.06-amd-amdhsa < %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=amdgpu9.00-amd-amdhsa < %t/on.ll | FileCheck --check-prefix=GFX900 %s
+
+; Target directives for supported target
+; SRAMECC-ON: .amdgcn_target "amdgpu9.06-amd-amdhsa-unknown-gfx906:sramecc+"
+; SRAMECC-OFF: .amdgcn_target "amdgpu9.06-amd-amdhsa-unknown-gfx906:sramecc-"
+; SRAMECC-ANY: .amdgcn_target "amdgpu9.06-amd-amdhsa-unknown-gfx906"
+
+; Unsupported target ignores the flag
+; GFX900: .amdgcn_target "amdgpu9.00-amd-amdhsa-unknown-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-no-on-off-modes.ll b/llvm/test/CodeGen/AMDGPU/module-flag-xnack-no-on-off-modes.ll
new file mode 100644
index 0000000000000..ba6cc11e6e2c0
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/module-flag-xnack-no-on-off-modes.ll
@@ -0,0 +1,54 @@
+; Test targets without xnack on/off mode support ignore module flags
+; Targets with only FEATURE_XNACK (but not FEATURE_XNACK_ON_OFF_MODES)
+; have xnack always on and ignore module flag settings.
+; The target ID should not contain the xnack specifier.
+
+; RUN: split-file %s %t
+; RUN: llc -mtriple=amdgpu12.5-amd-amdhsa < %t/on.ll | FileCheck --check-prefix=CHECK %s
+; RUN: llc -mtriple=amdgpu12.5-amd-amdhsa < %t/off.ll | FileCheck --check-prefix=CHECK %s
+; RUN: llc -mtriple=amdgpu12.5-amd-amdhsa < %t/absent.ll | FileCheck --check-prefix=CHECK %s
+
+; RUN: llc -mtriple=amdgpu12.50-amd-amdhsa < %t/on.ll | FileCheck --check-prefix=CHECK %s
+; RUN: llc -mtriple=amdgpu12.50-amd-amdhsa < %t/off.ll | FileCheck --check-prefix=CHECK %s
+; RUN: llc -mtriple=amdgpu12.50-amd-amdhsa < %t/absent.ll | FileCheck --check-prefix=CHECK %s
+
+; RUN: llc -mtriple=amdgpu12.51-amd-amdhsa < %t/on.ll | FileCheck --check-prefix=CHECK %s
+; RUN: llc -mtriple=amdgpu12.51-amd-amdhsa < %t/off.ll | FileCheck --check-prefix=CHECK %s
+; RUN: llc -mtriple=amdgpu12.51-amd-amdhsa < %t/absent.ll | FileCheck --check-prefix=CHECK %s
+
+; Module flags are ignored - target ID has no xnack specifier
+; CHECK: .amdgcn_target "amdgpu12.5{{[0-1]?}}-amd-amdhsa-unknown-gfx{{12-5-generic|1250|1251}}"
+
+; Targets without on/off mode support always have xnack enabled,
+; so loads must not overwrite pointer arguments (same behavior regardless of module flag)
+; CHECK-LABEL: {{^}}simple_clause:
+; CHECK: flat_load_b32 v4, v[0:1]
+; CHECK-NEXT: flat_load_b32 v5, 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/module-flag-xnack-sramecc-combined.ll b/llvm/test/CodeGen/AMDGPU/module-flag-xnack-sramecc-combined.ll
new file mode 100644
index 0000000000000..bd9130955e6d5
--- /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-unknown-gfx90a:sramecc-:xnack+"
+; CHECK: amdhsa.target: 'amdgcn-amd-amdhsa-unknown-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..b834e7714274a
--- /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-unknown-gfx900:xnack+"
+; XNACK-OFF: .amdgcn_target "amdgcn-amd-amdhsa-unknown-gfx900:xnack-"
+; XNACK-ANY: .amdgcn_target "amdgcn-amd-amdhsa-unknown-gfx900"
+
+; GFX801: .amdgcn_target "amdgcn-amd-amdhsa-unknown-gfx801:xnack+"
+; GFX803: .amdgcn_target "amdgcn-amd-amdhsa-unknown-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 f309c97d4b312..87004befd0d9e 100644
--- a/llvm/test/CodeGen/AMDGPU/nsa-reassign.ll
+++ b/llvm/test/CodeGen/AMDGPU/nsa-reassign.ll
@@ -1,4 +1,4 @@
-; RUN: llc -mtriple=amdgpu10.10 -mattr=-xnack -enable-misched=0 < %s | FileCheck -check-prefix=GCN %s
+; RUN: llc -mtriple=amdgpu10.10 --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 c5a5c2b7ede34..e63ee011caadc 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=amdgpu10.10 -mattr=-xnack -verify-machineinstrs -run-pass post-RA-hazard-rec -o - %s | FileCheck -check-prefix=GCN %s
+# RUN: llc -mtriple=amdgpu10.10 --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 e0d1a105aa29a..5c2c3b1722d33 100644
--- a/llvm/test/CodeGen/AMDGPU/occupancy-levels.ll
+++ b/llvm/test/CodeGen/AMDGPU/occupancy-levels.ll
@@ -1,12 +1,12 @@
-; RUN: llc -mtriple=amdgpu9.00 -mattr=-xnack < %s | FileCheck --check-prefixes=GCN,GFX9 %s
-; RUN: llc -mtriple=amdgpu9.50 -mattr=-xnack < %s | FileCheck --check-prefixes=GCN,GFX950 %s
+; RUN: llc -mtriple=amdgpu9.00 --amdgpu-xnack=false < %s | FileCheck --check-prefixes=GCN,GFX9 %s
+; RUN: llc -mtriple=amdgpu9.50 --amdgpu-xnack=false < %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=amdgpu9.00-amd-amdhsa -mattr=-xnack < %s | FileCheck --check-prefixes=GCN,GFX9TRAP %s
-; RUN: llc -mtriple=amdgpu9.50-amd-amdhsa -mattr=-xnack < %s | FileCheck --check-prefixes=GCN,GFX950TRAP %s
-; RUN: llc -mtriple=amdgpu10.10 -mattr=-xnack < %s | FileCheck --check-prefixes=GCN,GFX10,GFX10W32,GFX1010,GFX1010W32 %s
-; RUN: llc -mtriple=amdgpu10.10 -mattr=-xnack -mattr=+wavefrontsize64 < %s | FileCheck --check-prefixes=GCN,GFX10,GFX10W64,GFX1010,GFX1010W64 %s
+; RUN: llc -mtriple=amdgpu9.00-amd-amdhsa --amdgpu-xnack=false < %s | FileCheck --check-prefixes=GCN,GFX9TRAP %s
+; RUN: llc -mtriple=amdgpu9.50-amd-amdhsa --amdgpu-xnack=false < %s | FileCheck --check-prefixes=GCN,GFX950TRAP %s
+; RUN: llc -mtriple=amdgpu10.10 --amdgpu-xnack=false < %s | FileCheck --check-prefixes=GCN,GFX10,GFX10W32,GFX1010,GFX1010W32 %s
+; RUN: llc -mtriple=amdgpu10.10 --amdgpu-xnack=false -mattr=+wavefrontsize64 < %s | FileCheck --check-prefixes=GCN,GFX10,GFX10W64,GFX1010,GFX1010W64 %s
; RUN: llc -mtriple=amdgpu10.30 < %s | FileCheck --check-prefixes=GCN,GFX10,GFX10W32,GFX1030,GFX1030W32 %s
; RUN: llc -mtriple=amdgpu10.30 -mattr=+wavefrontsize64 < %s | FileCheck --check-prefixes=GCN,GFX10,GFX10W64,GFX1030,GFX1030W64 %s
; RUN: llc -mtriple=amdgpu11.00 < %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 2f2054afe010c..0ef9fc8f60351 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=amdgpu9.00 -mattr=+xnack -amdgpu-max-memory-clause=0 < %s | FileCheck -enable-var-scope -check-prefix=GCN %s
+; RUN: llc -mtriple=amdgpu9.00 --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 40c9500d54906..3f6dc826ad6ac 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=amdgpu6.00--amdpal < %s | FileCheck -check-prefix=SI %s
-; RUN: llc -mtriple=amdgpu8.02--amdpal -mattr=-flat-for-global,-xnack < %s | FileCheck -check-prefix=SI %s
+; RUN: llc -mtriple=amdgpu8.02--amdpal -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
diff erences.
diff --git a/llvm/test/CodeGen/AMDGPU/s_mulk_i32.ll b/llvm/test/CodeGen/AMDGPU/s_mulk_i32.ll
index 53681a9724881..d1ade036f1934 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=amdgpu6.00--amdpal < %s | FileCheck -check-prefix=GFX6 %s
-; RUN: llc -mtriple=amdgpu8.02--amdpal -mattr=-flat-for-global,-xnack < %s | FileCheck -check-prefix=GFX8 %s
+; RUN: llc -mtriple=amdgpu8.02--amdpal -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 e450aa86da0c0..e9bb77f3db5dd 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=amdgpu9.00-amd-amdhsa -mattr=+xnack -amdgpu-use-amdgpu-trackers=1 2>&1 < %s | FileCheck -check-prefixes=ERR-GCNTRACKERS %s
-; RUN: llc -mtriple=amdgpu9.00-amd-amdhsa -mattr=+xnack 2>&1 < %s | FileCheck -check-prefixes=GCN %s
+; RUN: not llc -mtriple=amdgpu9.00-amd-amdhsa --amdgpu-xnack=true -amdgpu-use-amdgpu-trackers=1 2>&1 < %s | FileCheck -check-prefixes=ERR-GCNTRACKERS %s
+; RUN: llc -mtriple=amdgpu9.00-amd-amdhsa --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 7c35952c36d3c..13c17b415f6d7 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=amdgpu9.06-amd-amdhsa -mattr=+xnack -run-pass=si-form-memory-clauses -verify-machineinstrs -o - %s | FileCheck %s
-# RUN: llc -mtriple=amdgpu9.06-amd-amdhsa -mattr=+xnack -passes="si-form-memory-clauses" -o - %s | FileCheck %s
+# RUN: llc -mtriple=amdgpu9.06-amd-amdhsa --amdgpu-xnack=true -run-pass=si-form-memory-clauses -verify-machineinstrs -o - %s | FileCheck %s
+# RUN: llc -mtriple=amdgpu9.06-amd-amdhsa --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 9c119f6ad326c..5c829d7ec88ef 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=amdgpu6.01 -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=amdgpu8.02 -enable-misched=0 -post-RA-scheduler=0 -amdgpu-spill-sgpr-to-vgpr=0 < %s | FileCheck --check-prefix=CHECK %s
-; RUN: llc -mtriple=amdgpu9.00 -mattr=-xnack,+enable-flat-scratch -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=amdgpu9.00 --amdgpu-xnack=false -mattr=+enable-flat-scratch -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=amdgpu10.30 -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 cbda67730e5a5..318cece479b1b 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=amdgpu9.00 < %s | FileCheck -check-prefixes=GCN,NO-ECC %s
-; RUN: llc -mtriple=amdgpu9.00 -mattr=+sramecc < %s | FileCheck -check-prefixes=GCN,NO-ECC %s
-; RUN: llc -mtriple=amdgpu9.00 -mattr=-sramecc < %s | FileCheck -check-prefixes=GCN,NO-ECC %s
-; RUN: llc -mtriple=amdgpu9.02 -mattr=+sramecc < %s | FileCheck -check-prefixes=GCN,NO-ECC %s
-; RUN: llc -mtriple=amdgpu9.04 -mattr=+sramecc < %s | FileCheck -check-prefixes=GCN,NO-ECC %s
-; RUN: llc -mtriple=amdgpu9.06 -mattr=+sramecc < %s | FileCheck -check-prefixes=GCN,ECC %s
-; RUN: llc -mtriple=amdgpu9.06 -mattr=-sramecc < %s | FileCheck -check-prefixes=GCN,NO-ECC %s
+; RUN: llc -mtriple=amdgpu9.00 -amdgpu-sramecc=1 < %s | FileCheck -check-prefixes=GCN,NO-ECC %s
+; RUN: llc -mtriple=amdgpu9.00 -amdgpu-sramecc=0 < %s | FileCheck -check-prefixes=GCN,NO-ECC %s
+
+; RUN: llc -mtriple=amdgpu9.06 -amdgpu-sramecc=1 < %s | FileCheck -check-prefixes=GCN,ECC %s
+; RUN: llc -mtriple=amdgpu9.06 -amdgpu-sramecc=0 < %s | FileCheck -check-prefixes=GCN,NO-ECC %s
; RUN: llc -mtriple=amdgpu12.50 < %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 56b591d2bf57f..0c6ba648bd9cf 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=amdgpu7.00 -debug-only=gcn-subtarget -o - %s 2>&1 | FileCheck --check-prefix=WARN %s
; RUN: llc -mtriple=amdgpu9.06 -debug-only=gcn-subtarget -o - %s 2>&1 | FileCheck --check-prefix=OFF %s
; RUN: llc -mtriple=amdgpu9.08 -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 5d9f57e9fecf8..a22a3ad48f43a 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=amdgpu7.00 -debug-only=gcn-subtarget -o - %s 2>&1 | FileCheck --check-prefix=WARN %s
; RUN: llc -mtriple=amdgpu9.06 -debug-only=gcn-subtarget -o - %s 2>&1 | FileCheck --check-prefix=ON %s
; RUN: llc -mtriple=amdgpu9.08 -debug-only=gcn-subtarget -o - %s 2>&1 | FileCheck --check-prefix=ON %s
; RUN: llc -mtriple=amdgpu12.50 -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 92526959bdb16..d4680f6e825b2 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=amdgpu12.51-amd-amdhsa < %s | FileCheck --check-prefix=GFX1251 %s
; RUN: llc -mtriple=amdgpu12.5-amd-amdhsa < %s | FileCheck --check-prefix=GFX125GEN %s
-; Even with -mattr=+xnack or -mattr=-xnack, the target ID doesn't change
-; RUN: llc -mtriple=amdgpu12.50-amd-amdhsa -mattr=+xnack < %s | FileCheck --check-prefix=GFX1250 %s
-; RUN: llc -mtriple=amdgpu12.50-amd-amdhsa -mattr=-xnack < %s | FileCheck --check-prefix=GFX1250 %s
-; RUN: llc -mtriple=amdgpu12.51-amd-amdhsa -mattr=+xnack < %s | FileCheck --check-prefix=GFX1251 %s
-; RUN: llc -mtriple=amdgpu12.51-amd-amdhsa -mattr=-xnack < %s | FileCheck --check-prefix=GFX1251 %s
-; RUN: llc -mtriple=amdgpu12.5-amd-amdhsa -mattr=+xnack < %s | FileCheck --check-prefix=GFX125GEN %s
-; RUN: llc -mtriple=amdgpu12.5-amd-amdhsa -mattr=-xnack < %s | FileCheck --check-prefix=GFX125GEN %s
+; Even with --amdgpu-xnack=true or --amdgpu-xnack=false, the target ID doesn't change
+; RUN: llc -mtriple=amdgpu12.50-amd-amdhsa --amdgpu-xnack=true < %s | FileCheck --check-prefix=GFX1250 %s
+; RUN: llc -mtriple=amdgpu12.50-amd-amdhsa --amdgpu-xnack=false < %s | FileCheck --check-prefix=GFX1250 %s
+; RUN: llc -mtriple=amdgpu12.51-amd-amdhsa --amdgpu-xnack=true < %s | FileCheck --check-prefix=GFX1251 %s
+; RUN: llc -mtriple=amdgpu12.51-amd-amdhsa --amdgpu-xnack=false < %s | FileCheck --check-prefix=GFX1251 %s
+; RUN: llc -mtriple=amdgpu12.5-amd-amdhsa --amdgpu-xnack=true < %s | FileCheck --check-prefix=GFX125GEN %s
+; RUN: llc -mtriple=amdgpu12.5-amd-amdhsa --amdgpu-xnack=false < %s | FileCheck --check-prefix=GFX125GEN %s
; GFX1250: .amdgcn_target "amdgpu12.50-amd-amdhsa-unknown-gfx1250"
; GFX1251: .amdgcn_target "amdgpu12.51-amd-amdhsa-unknown-gfx1251"
diff --git a/llvm/test/CodeGen/AMDGPU/tid-kd-xnack-off.ll b/llvm/test/CodeGen/AMDGPU/tid-kd-xnack-off.ll
index 87533d3fcfcba..dac5da844032a 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=amdgpu9.0a-amd-amdhsa -mattr=-xnack < %s | FileCheck --check-prefixes=ASM %s
-; RUN: llc -mtriple=amdgpu9.0a-amd-amdhsa -mattr=-xnack --filetype=obj < %s | llvm-objdump -s -j .rodata - | FileCheck --check-prefixes=OBJ %s
-; RUN: llc -mtriple=amdgpu9.0a-amd-amdhsa -mattr=-xnack --filetype=obj < %s | llvm-readelf --notes - | FileCheck --check-prefixes=ELF %s
+; RUN: llc -mtriple=amdgpu9.0a-amd-amdhsa < %s | FileCheck --check-prefixes=ASM %s
+; RUN: llc -mtriple=amdgpu9.0a-amd-amdhsa --filetype=obj < %s | llvm-objdump -s -j .rodata - | FileCheck --check-prefixes=OBJ %s
+; RUN: llc -mtriple=amdgpu9.0a-amd-amdhsa --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 01be974310a1f..6ce5cd774ae08 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=amdgpu9.0a-amd-amdhsa -mattr=+xnack < %s | FileCheck --check-prefixes=ASM %s
-; RUN: llc -mtriple=amdgpu9.0a-amd-amdhsa -mattr=+xnack --filetype=obj < %s | llvm-objdump -s -j .rodata - | FileCheck --check-prefixes=OBJ %s
-; RUN: llc -mtriple=amdgpu9.0a-amd-amdhsa -mattr=+xnack --filetype=obj < %s | llvm-readelf --notes - | FileCheck --check-prefixes=ELF %s
+; RUN: llc -mtriple=amdgpu9.0a-amd-amdhsa --amdgpu-xnack=true < %s | FileCheck --check-prefixes=ASM %s
+; RUN: llc -mtriple=amdgpu9.0a-amd-amdhsa --amdgpu-xnack=true --filetype=obj < %s | llvm-objdump -s -j .rodata - | FileCheck --check-prefixes=OBJ %s
+; RUN: llc -mtriple=amdgpu9.0a-amd-amdhsa --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 162581887abee..afce88383e875 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 8e858dea9ac28..66e5c3207d26b 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 6aec588a9e7b6..127a436dcca77 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 0ad5276b5a32f..bd0a8703a4999 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 7278f7ddbb868..58209e48b47e9 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 6107d75fd1618..d0d0881dade4a 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 3f5d651e3f53b..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=amdgpu9.00-amd-amdhsa < %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 c81b8b4617343..0fe041b972593 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 d5e72562866ce..9b03b3001104f 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 cfb0a2649ea26..81f1e66db0a69 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=amdgpu6.00 -debug-only=gcn-subtarget -filetype=null %s 2>&1 | FileCheck --check-prefix=WARN %s
-; RUN: llc -mtriple=amdgpu7.00 -debug-only=gcn-subtarget -filetype=null %s 2>&1 | FileCheck --check-prefix=WARN %s
; RUN: llc -mtriple=amdgpu8.01 -debug-only=gcn-subtarget -o - %s 2>&1 | FileCheck --check-prefix=OFF %s
; RUN: llc -mtriple=amdgpu9.00 -debug-only=gcn-subtarget -o - %s 2>&1 | FileCheck --check-prefix=OFF %s
; RUN: llc -mtriple=amdgpu9.06 -debug-only=gcn-subtarget -o - %s 2>&1 | FileCheck --check-prefix=OFF %s
; RUN: llc -mtriple=amdgpu10.10 -debug-only=gcn-subtarget -o - %s 2>&1 | FileCheck --check-prefix=OFF %s
-; RUN: llc -mtriple=amdgpu11.00 -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 9a76aa0889973..2ee1325884cb2 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=amdgpu6.00 -debug-only=gcn-subtarget -filetype=null %s 2>&1 | FileCheck --check-prefix=WARN %s
-; RUN: llc -mtriple=amdgpu7.00 -debug-only=gcn-subtarget -filetype=null %s 2>&1 | FileCheck --check-prefix=WARN %s
; RUN: llc -mtriple=amdgpu8.01 -debug-only=gcn-subtarget -o - %s 2>&1 | FileCheck --check-prefix=ON %s
; RUN: llc -mtriple=amdgpu9.00 -debug-only=gcn-subtarget -o - %s 2>&1 | FileCheck --check-prefix=ON %s
; RUN: llc -mtriple=amdgpu9.06 -debug-only=gcn-subtarget -o - %s 2>&1 | FileCheck --check-prefix=ON %s
; RUN: llc -mtriple=amdgpu10.10 -debug-only=gcn-subtarget -o - %s 2>&1 | FileCheck --check-prefix=ON %s
-; RUN: llc -mtriple=amdgpu11.00 -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..8bc8911d9a905 100644
--- a/llvm/test/MC/AMDGPU/xnack-mask.s
+++ b/llvm/test/MC/AMDGPU/xnack-mask.s
@@ -1,10 +1,10 @@
-// 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=amdgpu6.00 %s -filetype=null 2>&1 | FileCheck -check-prefix=NOSICIVI10 --implicit-check-not=error: %s
+// RUN: not llvm-mc -triple=amdgpu7.01 %s -filetype=null 2>&1 | FileCheck -check-prefix=NOSICIVI10 --implicit-check-not=error: %s
+// RUN: not llvm-mc -triple=amdgpu8.02 %s -filetype=null 2>&1 | FileCheck -check-prefix=NOSICIVI10 --implicit-check-not=error: %s
+// RUN: not llvm-mc -triple=amdgpu10.10 -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
+// RUN: not llvm-mc -triple=amdgpu8.10 -mattr=+xnack %s -filetype=null 2>&1 | FileCheck -check-prefix=XNACKERR --implicit-check-not=error: %s
+// RUN: not llvm-mc -triple=amdgpu8.10 -mattr=+xnack -show-encoding %s | FileCheck -check-prefix=XNACK %s
s_mov_b64 xnack_mask, -1
// NOSICIVI10: :[[@LINE-1]]:{{[0-9]+}}: error: xnack_mask register not available on this GPU
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 cfe-commits
mailing list