llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang-driver Author: Matt Arsenault (arsenm) <details> <summary>Changes</summary> 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@<!-- -->anthropic.com> --- Patch is 162.61 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/204595.diff 98 Files Affected: - (modified) clang/include/clang/Basic/TargetOptions.h (+16) - (modified) clang/include/clang/Options/Options.td (+12-4) - (modified) clang/lib/Basic/Targets/AMDGPU.cpp (+10) - (modified) clang/lib/CodeGen/CodeGenModule.cpp (+21) - (modified) clang/lib/Driver/ToolChains/Clang.cpp (+13) - (modified) clang/lib/Driver/ToolChains/Clang.h (+2) - (modified) clang/lib/Frontend/CompilerInvocation.cpp (+27) - (modified) clang/test/CIR/CodeGenHIP/target-features.hip (+9-8) - (modified) clang/test/CodeGenCXX/dynamic-cast-address-space.cpp (+2-2) - (added) clang/test/CodeGenOpenCL/amdgpu-module-flag-xnack-sramecc.cl (+22) - (added) clang/test/CodeGenOpenCL/amdgpu-xnack-any-only.cl (+27) - (modified) clang/test/Driver/amdgpu-features.c (+4-4) - (modified) clang/test/Driver/amdgpu-openmp-toolchain.c (+1-1) - (modified) clang/test/Driver/amdgpu-toolchain.c (+2-2) - (modified) clang/test/Driver/amdgpu-xnack-sramecc-flags.c (+36-33) - (modified) clang/test/Driver/hip-sanitize-options.hip (+17-17) - (modified) clang/test/Driver/hip-target-id.hip (+6-8) - (modified) clang/test/Driver/hip-toolchain-features.hip (+6-10) - (modified) clang/test/Driver/target-id.cl (+7-3) - (modified) llvm/docs/AMDGPUUsage.rst (+39-1) - (modified) llvm/docs/ReleaseNotes.md (+2) - (modified) llvm/lib/IR/Verifier.cpp (+19) - (modified) llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp (+33-49) - (modified) llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp (+45-1) - (modified) llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.h (+5) - (modified) llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp (+1-1) - (modified) llvm/lib/Target/AMDGPU/GCNSubtarget.cpp (+20-7) - (modified) llvm/lib/Target/AMDGPU/GCNSubtarget.h (+5-1) - (modified) llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.cpp (+8) - (modified) llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.h (+2-4) - (modified) llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp (+47-4) - (modified) llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h (+13) - (modified) llvm/test/CodeGen/AMDGPU/GlobalISel/extractelement-stack-lower.ll (+1-1) - (modified) llvm/test/CodeGen/AMDGPU/amdpal-callable.ll (+3-3) - (modified) llvm/test/CodeGen/AMDGPU/break-smem-soft-clauses.mir (+2-2) - (modified) llvm/test/CodeGen/AMDGPU/callee-special-input-vgprs-packed.ll (+1-1) - (modified) llvm/test/CodeGen/AMDGPU/cluster-flat-loads-postra.mir (+2-2) - (modified) llvm/test/CodeGen/AMDGPU/cluster_stores.ll (+1-1) - (modified) llvm/test/CodeGen/AMDGPU/directive-amdgcn-target.ll (+52-52) - (modified) llvm/test/CodeGen/AMDGPU/elf-header-flags-sramecc.ll (+8-8) - (modified) llvm/test/CodeGen/AMDGPU/elf-header-flags-xnack.ll (+2-2) - (modified) llvm/test/CodeGen/AMDGPU/flat-saddr-load.ll (+2-2) - (modified) llvm/test/CodeGen/AMDGPU/flat-scratch-reg.ll (+11-11) - (modified) llvm/test/CodeGen/AMDGPU/gfx902-without-xnack.ll (+3-2) - (modified) llvm/test/CodeGen/AMDGPU/greedy-reverse-local-assignment.ll (+2-2) - (modified) llvm/test/CodeGen/AMDGPU/hazard-hidden-bundle.mir (+3-3) - (modified) llvm/test/CodeGen/AMDGPU/hazard-in-bundle.mir (+3-3) - (modified) llvm/test/CodeGen/AMDGPU/hsa-metadata-kernel-code-props.ll (+3-3) - (modified) llvm/test/CodeGen/AMDGPU/hsa-metadata-resource-usage-function-ordering.ll (+3-3) - (modified) llvm/test/CodeGen/AMDGPU/hsa-note-no-func.ll (+4-4) - (modified) llvm/test/CodeGen/AMDGPU/immv216.ll (+3-3) - (modified) llvm/test/CodeGen/AMDGPU/limit-soft-clause-reg-pressure.mir (+2-2) - (modified) llvm/test/CodeGen/AMDGPU/materialize-frame-index-sgpr.ll (+3-3) - (added) llvm/test/CodeGen/AMDGPU/mattr-xnack-sramecc-legacy.ll (+23) - (added) llvm/test/CodeGen/AMDGPU/module-flag-sramecc.ll (+66) - (added) llvm/test/CodeGen/AMDGPU/module-flag-xnack-sramecc-combined.ll (+15) - (added) llvm/test/CodeGen/AMDGPU/module-flag-xnack.ll (+75) - (modified) llvm/test/CodeGen/AMDGPU/nsa-reassign.ll (+1-1) - (modified) llvm/test/CodeGen/AMDGPU/nsa-vmem-hazard.mir (+1-1) - (modified) llvm/test/CodeGen/AMDGPU/occupancy-levels.ll (+4-9) - (modified) llvm/test/CodeGen/AMDGPU/post-ra-soft-clause-dbg-info.ll (+1-1) - (modified) llvm/test/CodeGen/AMDGPU/s_addk_i32.ll (+1-1) - (modified) llvm/test/CodeGen/AMDGPU/s_mulk_i32.ll (+1-1) - (modified) llvm/test/CodeGen/AMDGPU/schedule-amdgpu-tracker-physreg-crash.ll (+2-2) - (modified) llvm/test/CodeGen/AMDGPU/soft-clause-dbg-value.mir (+2-2) - (modified) llvm/test/CodeGen/AMDGPU/spill-scavenge-offset.ll (+1-1) - (modified) llvm/test/CodeGen/AMDGPU/sram-ecc-default.ll (+6-6) - (modified) llvm/test/CodeGen/AMDGPU/sramecc-subtarget-feature-disabled.ll (+3-2) - (modified) llvm/test/CodeGen/AMDGPU/sramecc-subtarget-feature-enabled.ll (+3-2) - (modified) llvm/test/CodeGen/AMDGPU/target-id-xnack-always-on.ll (+7-7) - (modified) llvm/test/CodeGen/AMDGPU/tid-kd-xnack-off.ll (+5-4) - (modified) llvm/test/CodeGen/AMDGPU/tid-kd-xnack-on.ll (+3-3) - (modified) llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-all-off.ll (+2-1) - (modified) llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-all-on.ll (+2-1) - (modified) llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-any-off-1.ll (+2-1) - (modified) llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-any-off-2.ll (+2-1) - (modified) llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-any-on-1.ll (+2-1) - (modified) llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-any-on-2.ll (+2-1) - (removed) llvm/test/CodeGen/AMDGPU/tid-mul-func-xnack-invalid-any-off-on.ll (-24) - (modified) llvm/test/CodeGen/AMDGPU/tid-one-func-xnack-off.ll (+2-1) - (modified) llvm/test/CodeGen/AMDGPU/tid-one-func-xnack-on.ll (+2-1) - (modified) llvm/test/CodeGen/AMDGPU/xnack-subtarget-feature-disabled.ll (+3-4) - (modified) llvm/test/CodeGen/AMDGPU/xnack-subtarget-feature-enabled.ll (+3-4) - (added) llvm/test/Linker/Inputs/amdgpu-sramecc-module-flag-0.ll (+6) - (added) llvm/test/Linker/Inputs/amdgpu-sramecc-module-flag-1.ll (+6) - (added) llvm/test/Linker/Inputs/amdgpu-sramecc-module-flag-any.ll (+6) - (added) llvm/test/Linker/Inputs/amdgpu-xnack-module-flag-0.ll (+6) - (added) llvm/test/Linker/Inputs/amdgpu-xnack-module-flag-1.ll (+6) - (added) llvm/test/Linker/Inputs/amdgpu-xnack-module-flag-any.ll (+6) - (added) llvm/test/Linker/amdgpu-sramecc-module-flag-0.ll (+24) - (added) llvm/test/Linker/amdgpu-sramecc-module-flag-1.ll (+23) - (added) llvm/test/Linker/amdgpu-sramecc-module-flag-any.ll (+25) - (added) llvm/test/Linker/amdgpu-xnack-module-flag-0.ll (+24) - (added) llvm/test/Linker/amdgpu-xnack-module-flag-1.ll (+23) - (added) llvm/test/Linker/amdgpu-xnack-module-flag-any.ll (+25) - (modified) llvm/test/MC/AMDGPU/xnack-mask.s (+1-1) - (added) llvm/test/Verifier/AMDGPU/module-flag-sramecc.ll (+46) - (added) llvm/test/Verifier/AMDGPU/module-flag-xnack.ll (+46) ``````````diff diff --git a/clang/include/clang/Basic/TargetOptions.h b/clang/include/clang/Basic/TargetOptions.h index 51a2cfe14976c..466f0ef22f265 100644 --- a/clang/include/clang/Basic/TargetOptions.h +++ b/clang/include/clang/Basic/TargetOptions.h @@ -92,6 +92,22 @@ class TargetOptions { /// \brief AMDGPU Printf lowering scheme AMDGPUPrintfKind AMDGPUPrintfKindVal = AMDGPUPrintfKind::Hostcall; + /// \brief Enumeration values for AMDGPU xnack/sramecc settings + enum class AMDGPUFeatureState { + /// Feature state not specified and should generate most compatible code. + Any = 0, + /// Feature explicitly disabled + Disabled = 1, + /// Feature explicitly enabled + Enabled = 2 + }; + + /// \brief AMDGPU xnack setting from -mxnack/-mno-xnack + AMDGPUFeatureState AMDGPUXnackState = AMDGPUFeatureState::Any; + + /// \brief AMDGPU sramecc setting from -msramecc/-mno-sramecc + AMDGPUFeatureState AMDGPUSramEccState = AMDGPUFeatureState::Any; + // The code model to be used as specified by the user. Corresponds to // CodeModel::Model enum defined in include/llvm/Support/CodeGen.h, plus // "default" for the case when the user has not explicitly specified a diff --git a/clang/include/clang/Options/Options.td b/clang/include/clang/Options/Options.td index 0a24f8b35a3ea..7e2e69cc33582 100644 --- a/clang/include/clang/Options/Options.td +++ b/clang/include/clang/Options/Options.td @@ -5879,10 +5879,18 @@ defm cumode : SimpleMFlag<"cumode", " execution mode (AMDGPU only)", m_amdgpu_Features_Group>; defm tgsplit : SimpleMFlag<"tgsplit", "Enable", "Disable", " threadgroup split execution mode (AMDGPU only)", m_amdgpu_Features_Group>; -defm xnack : SimpleMFlag<"xnack", "Enable", "Disable", - " XNACK (AMDGPU only)", m_amdgpu_Features_Group>; -defm sramecc : SimpleMFlag<"sramecc", "Enable", "Disable", - " SRAMECC (AMDGPU only)", m_amdgpu_Features_Group>; +def mxnack : Flag<["-"], "mxnack">, Group<m_Group>, + Visibility<[ClangOption, CC1Option]>, + HelpText<"Enable XNACK (AMDGPU only)">; +def mno_xnack : Flag<["-"], "mno-xnack">, Group<m_Group>, + Visibility<[ClangOption, CC1Option]>, + HelpText<"Disable XNACK (AMDGPU only)">; +def msramecc : Flag<["-"], "msramecc">, Group<m_Group>, + Visibility<[ClangOption, CC1Option]>, + HelpText<"Enable SRAMECC (AMDGPU only)">; +def mno_sramecc : Flag<["-"], "mno-sramecc">, Group<m_Group>, + Visibility<[ClangOption, CC1Option]>, + HelpText<"Disable SRAMECC (AMDGPU only)">; defm wavefrontsize64 : SimpleMFlag<"wavefrontsize64", "Specify wavefront size 64", "Specify wavefront size 32", " mode (AMDGPU only)">; diff --git a/clang/lib/Basic/Targets/AMDGPU.cpp b/clang/lib/Basic/Targets/AMDGPU.cpp index bfa956fa9a4e3..6a87a67a688f9 100644 --- a/clang/lib/Basic/Targets/AMDGPU.cpp +++ b/clang/lib/Basic/Targets/AMDGPU.cpp @@ -229,6 +229,16 @@ AMDGPUTargetInfo::AMDGPUTargetInfo(const llvm::Triple &Triple, ReadOnlyFeatures.insert(F); } HalfArgsAndReturns = true; + + if (Opts.AMDGPUXnackState != TargetOptions::AMDGPUFeatureState::Any) { + OffloadArchFeatures["xnack"] = + Opts.AMDGPUXnackState == TargetOptions::AMDGPUFeatureState::Enabled; + } + + if (Opts.AMDGPUSramEccState != TargetOptions::AMDGPUFeatureState::Any) { + OffloadArchFeatures["sramecc"] = + Opts.AMDGPUSramEccState == TargetOptions::AMDGPUFeatureState::Enabled; + } } void AMDGPUTargetInfo::adjust(DiagnosticsEngine &Diags, LangOptions &Opts, diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index ddbebb1d70336..e095716a21f7e 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -1163,6 +1163,27 @@ void CodeGenModule::Release() { getModule().addModuleFlag(llvm::Module::Error, "amdgpu_printf_kind", MDStr); } + + const TargetOptions &TargetOpts = getTarget().getTargetOpts(); + + if (TargetOpts.AMDGPUXnackState != TargetOptions::AMDGPUFeatureState::Any) { + // TODO: Avoid emitting the xnack flag on targets which do not support + // xnack configuration. + getModule().addModuleFlag( + llvm::Module::Error, "amdgpu.xnack", + llvm::ConstantInt::get( + Int32Ty, TargetOpts.AMDGPUXnackState == + TargetOptions::AMDGPUFeatureState::Enabled)); + } + + if (TargetOpts.AMDGPUSramEccState != + TargetOptions::AMDGPUFeatureState::Any) { + getModule().addModuleFlag( + llvm::Module::Error, "amdgpu.sramecc", + llvm::ConstantInt::get( + Int32Ty, TargetOpts.AMDGPUSramEccState == + TargetOptions::AMDGPUFeatureState::Enabled)); + } } // Emit a global array containing all external kernels or device variables diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index e3288c81d4c95..d58b077c85acd 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -1535,6 +1535,15 @@ void Clang::AddARMTargetArgs(const llvm::Triple &Triple, const ArgList &Args, AddUnalignedAccessWarning(CmdArgs); } +void Clang::AddAMDGPUTargetArgs(const ArgList &Args, + ArgStringList &CmdArgs) const { + // Pass through -mxnack/-mno-xnack and -msramecc/-mno-sramecc flags to cc1. + if (Arg *A = Args.getLastArg(options::OPT_mxnack, options::OPT_mno_xnack)) + A->render(Args, CmdArgs); + if (Arg *A = Args.getLastArg(options::OPT_msramecc, options::OPT_mno_sramecc)) + A->render(Args, CmdArgs); +} + void Clang::RenderTargetOptions(const llvm::Triple &EffectiveTriple, const ArgList &Args, bool KernelOrKext, ArgStringList &CmdArgs) const { @@ -1562,6 +1571,10 @@ void Clang::RenderTargetOptions(const llvm::Triple &EffectiveTriple, AddAArch64TargetArgs(Args, CmdArgs); break; + case llvm::Triple::amdgcn: + AddAMDGPUTargetArgs(Args, CmdArgs); + break; + case llvm::Triple::loongarch32: case llvm::Triple::loongarch64: AddLoongArchTargetArgs(Args, CmdArgs); diff --git a/clang/lib/Driver/ToolChains/Clang.h b/clang/lib/Driver/ToolChains/Clang.h index 9adad5c5430f2..23f270ffb76a3 100644 --- a/clang/lib/Driver/ToolChains/Clang.h +++ b/clang/lib/Driver/ToolChains/Clang.h @@ -51,6 +51,8 @@ class LLVM_LIBRARY_VISIBILITY Clang : public Tool { void AddAArch64TargetArgs(const llvm::opt::ArgList &Args, llvm::opt::ArgStringList &CmdArgs) const; + void AddAMDGPUTargetArgs(const llvm::opt::ArgList &Args, + llvm::opt::ArgStringList &CmdArgs) const; void AddARMTargetArgs(const llvm::Triple &Triple, const llvm::opt::ArgList &Args, llvm::opt::ArgStringList &CmdArgs, diff --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp index d2847739e3143..e8f9bc5a60737 100644 --- a/clang/lib/Frontend/CompilerInvocation.cpp +++ b/clang/lib/Frontend/CompilerInvocation.cpp @@ -5024,6 +5024,18 @@ static void GenerateTargetArgs(const TargetOptions &Opts, if (!Opts.DarwinTargetVariantSDKVersion.empty()) GenerateArg(Consumer, OPT_darwin_target_variant_sdk_version_EQ, Opts.DarwinTargetVariantSDKVersion.getAsString()); + + // Generate AMDGPU xnack and sramecc flags. + if (Opts.AMDGPUXnackState == TargetOptions::AMDGPUFeatureState::Enabled) + GenerateArg(Consumer, OPT_mxnack); + else if (Opts.AMDGPUXnackState == TargetOptions::AMDGPUFeatureState::Disabled) + GenerateArg(Consumer, OPT_mno_xnack); + + if (Opts.AMDGPUSramEccState == TargetOptions::AMDGPUFeatureState::Enabled) + GenerateArg(Consumer, OPT_msramecc); + else if (Opts.AMDGPUSramEccState == + TargetOptions::AMDGPUFeatureState::Disabled) + GenerateArg(Consumer, OPT_mno_sramecc); } static bool ParseTargetArgs(TargetOptions &Opts, ArgList &Args, @@ -5055,6 +5067,21 @@ static bool ParseTargetArgs(TargetOptions &Opts, ArgList &Args, Opts.DarwinTargetVariantSDKVersion = Version; } + if (Arg *A = Args.getLastArg(options::OPT_mxnack, options::OPT_mno_xnack)) { + bool IsEnabled = A->getOption().matches(options::OPT_mxnack); + Opts.AMDGPUXnackState = IsEnabled + ? TargetOptions::AMDGPUFeatureState::Enabled + : TargetOptions::AMDGPUFeatureState::Disabled; + } + + if (Arg *A = + Args.getLastArg(options::OPT_msramecc, options::OPT_mno_sramecc)) { + bool IsEnabled = A->getOption().matches(options::OPT_msramecc); + Opts.AMDGPUSramEccState = IsEnabled + ? TargetOptions::AMDGPUFeatureState::Enabled + : TargetOptions::AMDGPUFeatureState::Disabled; + } + return Diags.getNumErrors() == NumErrorsBefore; } diff --git a/clang/test/CIR/CodeGenHIP/target-features.hip b/clang/test/CIR/CodeGenHIP/target-features.hip index 8d414edcd8e2c..afce90caca435 100644 --- a/clang/test/CIR/CodeGenHIP/target-features.hip +++ b/clang/test/CIR/CodeGenHIP/target-features.hip @@ -18,17 +18,17 @@ // only the delta (differing features) is emitted on cir.target-features. // RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fclangir \ -// RUN: -fcuda-is-device -target-cpu gfx900 -target-feature +xnack \ +// RUN: -fcuda-is-device -target-cpu gfx900 -mxnack \ // RUN: -emit-cir %s -o %t-delta.cir // RUN: FileCheck --check-prefix=CIR-DELTA %s --input-file=%t-delta.cir // RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fclangir \ -// RUN: -fcuda-is-device -target-cpu gfx900 -target-feature +xnack \ +// RUN: -fcuda-is-device -target-cpu gfx900 -mxnack \ // RUN: -emit-llvm %s -o %t-delta.cir.ll // RUN: FileCheck --check-prefix=LLVM-DELTA %s --input-file=%t-delta.cir.ll // RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip \ -// RUN: -fcuda-is-device -target-cpu gfx900 -target-feature +xnack \ +// RUN: -fcuda-is-device -target-cpu gfx900 -mxnack \ // RUN: -emit-llvm %s -o %t-delta.ll // RUN: FileCheck --check-prefix=LLVM-DELTA %s --input-file=%t-delta.ll @@ -49,17 +49,18 @@ __device__ void device_fn() {} // LLVM-DAG: attributes #[[K_ATTR]] = {{.*}}"target-cpu"="gfx900" // LLVM-DAG: attributes #[[D_ATTR]] = {{.*}}"target-cpu"="gfx900" -// AMDGPU with gfx900 + an extra +xnack feature: only the delta is emitted. +// AMDGPU with gfx900 + xnack enabled via -mxnack: emitted as module flag. // CIR-DELTA: cir.func{{.*}} @_Z6kernelv() // CIR-DELTA-SAME: "cir.target-cpu" = "gfx900" -// CIR-DELTA-SAME: "cir.target-features" = "+xnack" +// CIR-DELTA-NOT: cir.target-features // CIR-DELTA: cir.func{{.*}} @_Z9device_fnv() // CIR-DELTA-SAME: "cir.target-cpu" = "gfx900" -// CIR-DELTA-SAME: "cir.target-features" = "+xnack" +// CIR-DELTA-NOT: cir.target-features // LLVM-DELTA: define{{.*}} void @_Z6kernelv(){{.*}} #[[K_ATTR_D:[0-9]+]] // LLVM-DELTA: define{{.*}} void @_Z9device_fnv(){{.*}} #[[D_ATTR_D:[0-9]+]] -// LLVM-DELTA-DAG: attributes #[[K_ATTR_D]] = {{.*}}"target-cpu"="gfx900"{{.*}}"target-features"="+xnack" -// LLVM-DELTA-DAG: attributes #[[D_ATTR_D]] = {{.*}}"target-cpu"="gfx900"{{.*}}"target-features"="+xnack" +// LLVM-DELTA-DAG: attributes #[[K_ATTR_D]] = {{.*}}"target-cpu"="gfx900" +// LLVM-DELTA-DAG: attributes #[[D_ATTR_D]] = {{.*}}"target-cpu"="gfx900" +// LLVM-DELTA-DAG: !{i32 1, !"amdgpu.xnack", i32 1} diff --git a/clang/test/CodeGenCXX/dynamic-cast-address-space.cpp b/clang/test/CodeGenCXX/dynamic-cast-address-space.cpp index e5873c1a8de90..2d48708601ccf 100644 --- a/clang/test/CodeGenCXX/dynamic-cast-address-space.cpp +++ b/clang/test/CodeGenCXX/dynamic-cast-address-space.cpp @@ -107,9 +107,9 @@ const B& f(A *a) { // CHECK: attributes #[[ATTR3]] = { nounwind } // CHECK: attributes #[[ATTR4]] = { noreturn } //. -// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR0]] = { mustprogress noinline optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+16-bit-insts,+add-min-max-insts,+ashr-pk-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-fmin-fmax-global-f32,+atomic-fmin-fmax-global-f64,+atomic-global-pk-add-bf16-inst,+bf16-cvt-insts,+bf16-pk-insts,+bf16-trans-insts,+bf8-cvt-scale-insts,+bitop3-insts,+ci-insts,+clusters,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot13-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+f16bf16-to-fp6bf6-cvt-scale-insts,+f32-to-f16bf16-cvt-sr-insts,+flat-global-insts,+fp4-cvt-scale-insts,+fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-cvt-scale-insts,+fp8-insts,+fp8e5m3-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx1251-gemm-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gfx950-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mcast-load-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+permlane16-swap,+permlane32-swap,+pk-add-min-max-insts,+prng-inst,+qsad-insts,+s-memrealtime,+s-memtime-inst,+s-wakeup-barrier-inst,+sad-insts,+setprio-inc-wg-inst,+swmmac-gfx1200-insts,+swmmac-gfx1250-insts,+tanh-insts,+tensor-cvt-lut-insts,+transpose-load-f4f6-insts,+vmem-pref-insts,+vmem-to-lds-load-insts,+wavefrontsize32,+wavefrontsize64,+wmma-128b-insts,+wmma-256b-insts,+xf32-insts" } +// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR0]] = { mustprogress noinline optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+16-bit-insts,+add-min-max-insts,+ashr-pk-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-fmin-fmax-global-f32,+atomic-fmin-fmax-global-f64,+atomic-global-pk-add-bf16-inst,+bf16-cvt-insts,+bf16-pk-insts,+bf16-trans-insts,+bf8-cvt-scale-insts,+bitop3-insts,+ci-insts,+clusters,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot13-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+f16bf16-to-fp6bf6-cvt-scale-insts,+f32-to-f16bf16-cvt-sr-insts,+flat-global-insts,+fp4-cvt-scale-insts,+fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-cvt-scale-insts,+fp8-insts,+fp8e5m3-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx1251-gemm-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gfx950-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mcast-load-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+permlane16-swap,+permlane32-swap,+pk-add-min-max-insts,+prng-inst,+qsad-insts,+s-memrealtime,+s-memtime-inst,+s-wakeup-barrier-inst,+sad-insts,+setprio-inc-wg-inst,+swmmac-gfx1200-insts,+swmmac-gfx1250-insts,+tanh-insts,+tensor-cvt-lut-insts,+transpose-load-f4f6-insts,+vmem-pref-insts,+vmem-to-lds-load-insts,+wavefrontsize32,+wavefrontsize64,+wmma-128b-insts,+wmma-256b-insts,+xf32-insts,+xnack-any-only" } // WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR1:[0-9]+]] = { nounwind willreturn memory(read) } -// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR2:[0-9]+]] = { "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+16-bit-insts,+add-min-max-insts,+ashr-pk-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-fmin-fmax-global-f32,+atomic-fmin-fmax-global-f64,+atomic-global-pk-add-bf16-inst,+bf16-cvt-insts,+bf16-pk-insts,+bf16-trans-insts,+bf8-cvt-scale-insts,+bitop3-insts,+ci-insts,+clusters,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot13-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+f16bf16-to-fp6bf6-cvt-scale-insts,+f32-to-f16bf16-cvt-sr-insts,+flat-global-insts,+fp4-cvt-scale-insts,+fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-cvt-scale-insts,+fp8-insts,+fp8e5m3-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx1251-gemm-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gfx950-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mcast-load-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+permlane16-swap,+permlane32-swap,+pk-add-min-max-insts,+prng-inst,+qsad-insts,+s-memrealtime,+s-memtime-inst,+s-wakeup-barrier-inst,+sad-insts,+setprio-inc-wg-inst,+swmmac-gfx1200-insts,+swmmac-gfx1250-insts,+tanh-insts,+tensor-cvt-lut-insts,+transpose-load-f4f6-insts,+vmem-pref-insts,+vmem-to-lds-load-insts,+wavefrontsize32,+wavefrontsize64,+wmma-128b-insts,+wmma-256b-insts,+xf32-insts" } +// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR2:[0-9]+]] = { "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+16-bit-insts,+add-min-max-insts,+ashr-pk-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-fmin-fmax-global-f32,+atomic-fmin-fmax-global-f64,+atomic-global-pk-add-bf16-inst,+bf16-cvt-insts,+bf16-pk-insts,+bf16-trans-insts,+bf8-cvt-scale-insts,+bitop3-insts,+ci-insts,+clusters,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot13-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+f16bf16-to-fp6bf6-cvt-scale-insts,+f32-to-f16bf16-cvt-sr-insts,+flat-global-insts,+fp4-cvt-scale-insts,+fp6bf6-cvt-scale-insts,+fp8-conversion-insts,+fp8-cvt-scale-insts,+fp8-insts,+fp8e5m3-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx1251-gemm-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gfx940-insts,+gfx950-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mcast-load-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+permlane16-swap,+permlane32-swap,+pk-add-min-max-insts,+prng-inst,+qsad-insts,+s-memrealtime,+s-memtime-inst,+s-wakeup-barrier-inst,+sad-insts,+setprio-inc-wg-inst,+swmmac-gfx1200-insts,+swmmac-gfx1250-insts,+tanh-insts,+tensor-cvt-lut-insts,+transpose-load-f4f6-insts,+vmem-pref-insts,+vmem-to-lds-load-insts,+wavefrontsize32,+wavefrontsize64,+wmma-128b-insts,+wmma-256b-insts,+xf32-insts,+xnack-any-only" } // WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR3]] = { nounwind } // WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR4]] = { noreturn } //. diff --git a/clang/test/CodeGenOpenCL/amdgpu-module-flag-xnack-sramecc.cl b/clang/test/CodeGenOpenCL/amdgpu-module-flag-xnack-sramecc.cl new file mode 100644 index 0000000000000..439a815f5c42d --- /dev/null +++ b/clang/test/CodeGenOpenCL/amdgpu-module-flag-xnack-sramecc.cl @@ -0,0 +1,22 @@ +// Test that xnack and sramecc module flags are emitted based on -m flags +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx90a \ +// RUN: -mxnack -mno-sramecc \ +// RUN: -emit-llvm -o - %s | FileCheck %s --check-prefixes=XNACK-ON,SRAMECC-OFF + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx90a \ +// RUN: -mno-xnack -msramecc \ +// RUN: -emit-llvm -o - %s | FileCheck %s --check-prefixes=XNACK-OFF,SRAMECC-ON + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx90a \ +// RUN: -emit-llvm -o - %s | FileCheck %s --check-prefix=NO-FLAGS + +// XNACK-ON-DAG: !{i32 1, !"amdgpu.xnack", i32 1} +// XNACK-OFF-DAG: !{i32 1, !"amdgpu.xnack", i32 0} +// SRAMECC-ON-DAG: !{i32 1, !"amdgpu.sramecc", i32 1} +// SRAMECC-OFF-DAG: !{i32 1, !"amdgpu.sramecc", i32 0} + +// When no explicit xnack/sramecc feature is set, no module flags are emitted +// NO-FLAGS-NOT: !"amdgpu.xnack" +// NO-FLAGS-NOT: !"amdgpu.sramecc" + +__attribute__((device)) void test() {} diff --git a/clang/test/CodeGenOpenCL/amdgpu-xnack-any-only.cl b/clang/test/CodeGenOpenCL/amdgpu-xnack-any-only.cl new file mode 100644 index 0000000000000..709ca1acf4cb3 --- /dev/null +++ b/clang/test/CodeGenOpenCL/amdgpu-xnack-any-only.cl @@ -0,0 +1,27 @@ +// Test that xnack module flags are emitted for all targets, regardless of support. +// Targets without FEATURE_XNACK_ON_OFF_MODES (like gfx12-5-generic, gfx1250, gfx1251) +// will ignore the module flag during codegen, but it is still emitted by clang. +// TODO: In the future, clang ... [truncated] `````````` </details> https://github.com/llvm/llvm-project/pull/204595 _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
