[Mlir-commits] [mlir] [mlir][ROCDL] Add TargetInfo to replace Chipset, allow features queries (PR #223562)
Krzysztof Drewniak
llvmlistbot at llvm.org
Thu Sep 17 09:03:53 PDT 2026
https://github.com/krzysz00 updated https://github.com/llvm/llvm-project/pull/223562
>From 8dbdda498f1efef3d7dde659f00aec2e0af82fe3 Mon Sep 17 00:00:00 2001
From: Krzysztof Drewniak <Krzysztof.Drewniak at amd.com>
Date: Mon, 14 Sep 2026 21:29:06 +0000
Subject: [PATCH 1/2] [mlir][ROCDL] Add TargetInfo to replace Chipset, allow
features queries
Add a now ROCDL::TargetInfo struct that parses AMDGPU triples and
target names using the same logic that Clang and LLVM
use (TargetParser) and maintains the set of features available on a
given GPU.
This is an improvement over the old `amdgpu::Chipset` struct since
that was just a version number and often became stale compared to the
knowledge exposed by LLVM, such as gfx1170 having OCP FP8 support even
though other gfx11 chips don't have it.
This struct also allows for moving to new-style
triples (amdgpu9.42-amd-amdhsa vs amdgcn-amd-amdhsa--gfx942, for
example), which is an ongoing migration in other parts of the compiler
that this PR lets us follow.
It also enables compiling for generic targets, like `gfx11-generic`,
which can be run on all chips in a generation.
The TargetInfo tracks the set of features available on a chip so you
can `.has()` a feature.
Finally, it copies over the handling of `:xnack` and `:sramecc`
modifiers from Clang, where the frontend converts them from "features"
into module flags. Here this is done with the
`migrateArchFeaturesToModuleFlags()` method, which downstream passes
will need to call.
This PR doesn't migrate any users to the new system, it's been split
for revewiability.
AI disclosure: I steered this, Claude wrote the code, I tried to clean
up the docs.
---
.../mlir/Dialect/LLVMIR/ROCDLTargetInfo.h | 193 +++++++++
mlir/lib/Dialect/LLVMIR/CMakeLists.txt | 2 +
.../lib/Dialect/LLVMIR/IR/ROCDLTargetInfo.cpp | 242 +++++++++++
mlir/unittests/Dialect/LLVMIR/CMakeLists.txt | 2 +
.../Dialect/LLVMIR/ROCDLTargetInfoTest.cpp | 392 ++++++++++++++++++
5 files changed, 831 insertions(+)
create mode 100644 mlir/include/mlir/Dialect/LLVMIR/ROCDLTargetInfo.h
create mode 100644 mlir/lib/Dialect/LLVMIR/IR/ROCDLTargetInfo.cpp
create mode 100644 mlir/unittests/Dialect/LLVMIR/ROCDLTargetInfoTest.cpp
diff --git a/mlir/include/mlir/Dialect/LLVMIR/ROCDLTargetInfo.h b/mlir/include/mlir/Dialect/LLVMIR/ROCDLTargetInfo.h
new file mode 100644
index 0000000000000..68dd721c74ac1
--- /dev/null
+++ b/mlir/include/mlir/Dialect/LLVMIR/ROCDLTargetInfo.h
@@ -0,0 +1,193 @@
+//===- ROCDLTargetInfo.h - AMDGPU target description ------------*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+#ifndef MLIR_DIALECT_LLVMIR_ROCDLTARGETINFO_H_
+#define MLIR_DIALECT_LLVMIR_ROCDLTARGETINFO_H_
+
+#include "mlir/IR/Diagnostics.h"
+#include "mlir/Support/LLVM.h"
+#include "llvm/TargetParser/AMDGPUTargetParser.h"
+#include "llvm/TargetParser/Triple.h"
+#include <optional>
+
+namespace mlir::ROCDL {
+
+/// Describes the AMDGPU target a lowering is producing code for: the triple's
+/// subarch (which identifies the GPU) together with the resolved set of
+/// frontend-visible target features.
+///
+/// Lowerings should gate on features (`has(FEAT_...)`) rather than on ISA
+/// version arithmetic, and add features if necessary.
+class TargetInfo {
+public:
+ using Feature = ::llvm::AMDGPU::AMDGPUFeature;
+
+ /// Constructs an unknown target: no subarch, and every feature query answers
+ /// false.
+ TargetInfo() = default;
+
+ /// Resolves a target description.
+ ///
+ /// \p arch names the architecture the way Clang does, and accepts any of:
+ ///
+ /// - a full target ID, "<triple>-<processor>[:<feature><+|->]*", such as
+ /// "amdgcn-amd-amdhsa--gfx90a:sramecc+:xnack-" (what `rocminfo` prints
+ /// for a device's ISA) or "amdgpu9.0a-amd-amdhsa--gfx90a";
+ /// - a triple on its own, such as "amdgpu9.42-amd-amdhsa" or the legacy
+ /// subarch-less "amdgcn-amd-amdhsa";
+ /// - a processor on its own, with optional target-ID modifiers: "gfx942",
+ /// "gfx942:xnack+", "gfx9-4-generic".
+ ///
+ /// Only xnack and sramecc may be given as modifiers, and only on a processor
+ /// that supports them; this is the same grammar `clang::parseTargetID`
+ /// accepts, and it is validated by `llvm::AMDGPU::TargetID`.
+ ///
+ /// \p waveSize pins the wavefront size for targets that run at either, and
+ /// must be 0 (meaning the target's own default), 32, or 64.
+ ///
+ /// Diagnostics are emitted via `emitError`.
+ static FailureOr<TargetInfo>
+ get(StringRef arch, unsigned waveSize = 0,
+ function_ref<InFlightDiagnostic()> emitError = nullptr);
+
+ /// Parses \p arch into a target ID, accepting the spellings `get()`
+ /// documents, or returns nullopt if it names no valid target.
+ ///
+ /// Use only if you need to get the individual components of the target ID.
+ static std::optional<::llvm::AMDGPU::TargetID> parseTargetID(StringRef arch);
+
+ /// Returns whether the target has \p feature.
+ bool has(Feature feature) const { return featureBits.test(feature); }
+
+ /// Returns whether the target's fp8 conversions exist and use the OCP formats
+ /// (E4M3FN/E5M2) rather than the FNUZ ones.
+ bool hasOcpFp8() const {
+ return has(::llvm::AMDGPU::FEAT_OCP_FP8_CONVERSION_INSTS);
+ }
+
+ /// Returns whether the target has fp8 conversions that use the FNUZ formats
+ /// (E4M3FNUZ/E5M2FNUZ).
+ bool hasFnuzFp8() const {
+ return has(::llvm::AMDGPU::FEAT_FP8_CONVERSION_INSTS) && !hasOcpFp8();
+ }
+
+ /// Returns whether the target belongs to gfx generation \p major (9 for any
+ /// gfx9xx, 12 for any gfx12xx, ...).
+ ///
+ /// Prefer `has()` where a feature expresses the condition; this is used when
+ /// no feature exists and the property being checked is a function of the
+ /// major ISA generation (such as the details of buffer encoding).
+ bool isGeneration(unsigned major) const;
+
+ /// Returns the width in bits of the num_records field of the buffer resource
+ /// (V#), or nullopt for an unknown target.
+ std::optional<unsigned> getBufferResourceNumRecordsWidth() const;
+
+ /// Returns the maximum LDS in bytes a single workgroup can address, or
+ /// nullopt for an unknown target.
+ std::optional<unsigned> getMaxAddressableLocalMemorySize() const;
+
+ /// Returns the wavefront size, or nullopt for an unknown target. Targets that
+ /// support both sizes report 32 unless "+wavefrontsize64" was requested.
+ std::optional<unsigned> getWavefrontSize() const;
+
+ /// Returns whether the GPU can be configured for 32-lane or 64-lane
+ /// wavefronts.
+ bool supportsBothWavefrontSizes() const { return dualWavefrontSize; }
+
+ /// Returns the total number of SGPRs, or nullopt for an unknown target.
+ std::optional<unsigned> getTotalNumSGPRs() const;
+
+ /// Returns the number of SGPRs addressable by a kernel, or nullopt for an
+ /// unknown target. This is below getTotalNumSGPRs() where some are reserved.
+ std::optional<unsigned> getAddressableNumSGPRs() const;
+
+ /// Returns the SGPR allocation granularity in registers, or nullopt for an
+ /// unknown target.
+ std::optional<unsigned> getSGPRAllocGranule() const;
+
+ /// Returns the VGPR allocation granularity in registers, or nullopt for an
+ /// unknown target. This property is wavesize-dependent.
+ std::optional<unsigned> getVGPRAllocGranule() const;
+
+ /// Returns the number of LDS banks per compute unit, or nullopt for an
+ /// unknown target.
+ std::optional<unsigned> getLDSBankCount() const;
+
+ /// Returns the maximum number of waves per execution unit, ignoring any
+ /// limits a particular kernel imposes, or nullopt for an unknown target.
+ std::optional<unsigned> getMaxWavesPerEU() const;
+
+ /// Returns whether xnack is on, off, either, or unsupported on this target.
+ /// "Any" means the target supports both and no `:xnack+/-` modifier was used.
+ ::llvm::AMDGPU::TargetIDSetting getXnackSetting() const {
+ return xnackSetting;
+ }
+
+ /// Returns whether sramecc is on, off, either, or unsupported, as for
+ /// getXnackSetting().
+ ::llvm::AMDGPU::TargetIDSetting getSramEccSetting() const {
+ return sramEccSetting;
+ }
+
+ /// Records the xnack and sramecc settings this target's ID pinned onto the
+ /// module \p op, as the `rocdl.xnack` and `rocdl.sramecc` attributes that
+ /// translate to the `amdgpu.xnack` and `amdgpu.sramecc` module flags.
+ ///
+ /// These flags are given as `:{xnack,sramecc}` target-ID "modifiers",
+ /// since they used to be subtarget features, but now frontends (like us and
+ /// Clang) need to migrate them into module flags. This representation keeps
+ /// us compatible with Clang and the output of tools like `rocminfo`.
+ ///
+ /// If a particular modifier is not given, no attribute is set for it, putting
+ /// that value into its "any" state if it is controllable.
+ void migrateArchFeaturesToModuleFlags(Operation *op) const;
+
+ /// Returns the ISA version. For a generic target this is the floor of the
+ /// family it covers (gfx9-4-generic reports 9.4.0), so it must not be used to
+ /// decide whether an instruction is available.
+ ::llvm::AMDGPU::IsaVersion getIsaVersion() const;
+
+ ::llvm::Triple::SubArchType getSubArch() const { return subArch; }
+ ::llvm::AMDGPU::GPUKind getGPUKind() const { return kind; }
+
+ /// Returns the canonical GPU name ("gfx942", "gfx9-4-generic"), or "" if the
+ /// target is unknown.
+ StringRef getArchName() const;
+
+ /// Returns whether this is a "gfxN-generic" target, which carries only the
+ /// features common to every GPU it covers.
+ bool isGeneric() const;
+
+ /// Returns whether no GPU was identified, in which case every feature query
+ /// answers false.
+ bool isUnknown() const { return kind == ::llvm::AMDGPU::GK_NONE; }
+
+ const ::llvm::AMDGPU::AMDGPUFeatureBitset &getFeatures() const {
+ return featureBits;
+ }
+
+private:
+ ::llvm::Triple::SubArchType subArch = ::llvm::Triple::NoSubArch;
+ ::llvm::AMDGPU::GPUKind kind = ::llvm::AMDGPU::GK_NONE;
+ ::llvm::AMDGPU::AMDGPUFeatureBitset featureBits;
+ ::llvm::AMDGPU::TargetIDSetting xnackSetting =
+ ::llvm::AMDGPU::TargetIDSetting::Unsupported;
+ ::llvm::AMDGPU::TargetIDSetting sramEccSetting =
+ ::llvm::AMDGPU::TargetIDSetting::Unsupported;
+ bool dualWavefrontSize = false;
+};
+
+/// Returns the target architecture that a pass should parse, given its `arch`
+/// option and the value of the deprecated alias that `arch` replaced.
+///
+/// The alias is only consulted when `arch` is left at "invalid".
+StringRef resolveArchOption(StringRef arch, StringRef deprecatedAlias);
+
+} // namespace mlir::ROCDL
+
+#endif // MLIR_DIALECT_LLVMIR_ROCDLTARGETINFO_H_
diff --git a/mlir/lib/Dialect/LLVMIR/CMakeLists.txt b/mlir/lib/Dialect/LLVMIR/CMakeLists.txt
index 50351fde8128f..6b7827d5ab12b 100644
--- a/mlir/lib/Dialect/LLVMIR/CMakeLists.txt
+++ b/mlir/lib/Dialect/LLVMIR/CMakeLists.txt
@@ -116,6 +116,7 @@ add_public_tablegen_target(MLIRROCDLOpsShardGen)
add_mlir_dialect_library(MLIRROCDLDialect
IR/ROCDLDialect.cpp
${ROCDL_SHARDED_SRCS}
+ IR/ROCDLTargetInfo.cpp
ADDITIONAL_HEADER_DIRS
${MLIR_MAIN_INCLUDE_DIR}/mlir/Dialect/LLVMIR
@@ -130,6 +131,7 @@ add_mlir_dialect_library(MLIRROCDLDialect
LINK_COMPONENTS
AsmParser
Core
+ TargetParser
LINK_LIBS PUBLIC
MLIRIR
diff --git a/mlir/lib/Dialect/LLVMIR/IR/ROCDLTargetInfo.cpp b/mlir/lib/Dialect/LLVMIR/IR/ROCDLTargetInfo.cpp
new file mode 100644
index 0000000000000..cb5cc47c3af13
--- /dev/null
+++ b/mlir/lib/Dialect/LLVMIR/IR/ROCDLTargetInfo.cpp
@@ -0,0 +1,242 @@
+//===- ROCDLTargetInfo.cpp - AMDGPU target description --------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "mlir/Dialect/LLVMIR/ROCDLTargetInfo.h"
+#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
+#include "mlir/Dialect/LLVMIR/ROCDLDialect.h"
+#include "mlir/IR/Builders.h"
+#include "llvm/ADT/SmallVector.h"
+#include "llvm/ADT/StringRef.h"
+#include "llvm/ADT/Twine.h"
+
+using namespace mlir;
+using namespace mlir::ROCDL;
+
+namespace AMDGPU = ::llvm::AMDGPU;
+using ::llvm::Triple;
+
+/// Reports \p message through \p emitError if it is non-null, and returns
+/// failure.
+static LogicalResult fail(function_ref<InFlightDiagnostic()> emitError,
+ const Twine &message) {
+ if (emitError)
+ emitError() << message;
+ return failure();
+}
+
+std::optional<AMDGPU::TargetID> TargetInfo::parseTargetID(StringRef arch) {
+ // If we see a five-component triple, that's maximally authoritative.
+ SmallVector<StringRef, 5> parts;
+ arch.split(parts, '-', /*MaxSplit=*/4);
+ if (parts.size() == 5)
+ return AMDGPU::TargetID::parseTargetIDString(arch);
+
+ // Otherwise, handle bare triples.
+ Triple triple(Triple::normalize(arch));
+ if (triple.isAMDGCN())
+ return AMDGPU::TargetID::parse(triple, "");
+
+ // Otherwise, take the implicit amdgcn-amd-amdhsa legacy triple and pair it
+ // with an arch name.
+ return AMDGPU::TargetID::parse(Triple("amdgcn-amd-amdhsa"), arch);
+}
+
+/// Pins the wavefront size in \p bits, mirroring the policy LLVM applies in
+/// fillAMDGCNFeatureMap: a target that only runs at one size rejects a request
+/// for the other, and a target that runs at either defaults to wave32.
+static LogicalResult
+resolveWavefrontSize(AMDGPU::AMDGPUFeatureBitset &bits, unsigned waveSize,
+ function_ref<InFlightDiagnostic()> emitError) {
+ bool targetWave32 = bits.test(AMDGPU::FEAT_WAVEFRONTSIZE32);
+ bool targetWave64 = bits.test(AMDGPU::FEAT_WAVEFRONTSIZE64);
+
+ switch (waveSize) {
+ case 0:
+ // A target that runs at either size and was not asked for one runs wave32.
+ if (!targetWave32 && !targetWave64)
+ bits.set(AMDGPU::FEAT_WAVEFRONTSIZE32);
+ return success();
+ case 32:
+ if (targetWave64)
+ return fail(emitError, "target only supports a wavefront size of 64");
+ bits.set(AMDGPU::FEAT_WAVEFRONTSIZE32);
+ return success();
+ case 64:
+ if (targetWave32)
+ return fail(emitError, "target only supports a wavefront size of 32");
+ bits.set(AMDGPU::FEAT_WAVEFRONTSIZE64);
+ return success();
+ default:
+ return fail(emitError,
+ "wavefront size must be 32 or 64, got " + Twine(waveSize));
+ }
+}
+
+FailureOr<TargetInfo>
+TargetInfo::get(StringRef arch, unsigned waveSize,
+ function_ref<InFlightDiagnostic()> emitError) {
+ if (arch.empty())
+ return fail(emitError, "target architecture cannot be empty");
+
+ std::optional<AMDGPU::TargetID> id = parseTargetID(arch);
+ if (!id)
+ return fail(emitError, "'" + arch +
+ "' is not a valid AMDGPU architecture: expected "
+ "a GPU name, a triple, or a target ID");
+
+ TargetInfo info;
+ info.kind = id->getGPUKind();
+ info.subArch = AMDGPU::getSubArch(info.kind);
+ info.featureBits = AMDGPU::getFeatureBitset(info.kind);
+ info.xnackSetting = id->getXnackSetting();
+ info.sramEccSetting = id->getSramEccSetting();
+
+ // Recorded before we record the user's choice since that lives in the same
+ // bitmap.
+ info.dualWavefrontSize =
+ !info.isUnknown() &&
+ !info.featureBits.test(AMDGPU::FEAT_WAVEFRONTSIZE32) &&
+ !info.featureBits.test(AMDGPU::FEAT_WAVEFRONTSIZE64);
+
+ if (info.isUnknown() && waveSize != 0 && waveSize != 32 && waveSize != 64)
+ return fail(emitError,
+ "wavefront size must be 32 or 64, got " + Twine(waveSize));
+ if (!info.isUnknown() &&
+ failed(resolveWavefrontSize(info.featureBits, waveSize, emitError)))
+ return failure();
+
+ return info;
+}
+
+StringRef mlir::ROCDL::resolveArchOption(StringRef arch,
+ StringRef deprecatedAlias) {
+ if (!arch.empty() && arch != "invalid")
+ return arch;
+ return deprecatedAlias.empty() ? arch : deprecatedAlias;
+}
+
+bool TargetInfo::isGeneration(unsigned major) const {
+ // The generation features are cumulative: a gfx12 target has every
+ // FEAT_GFX*_INSTS bit from gfx8 up to gfx12. So a target is *in* generation N
+ // when it has N's bit but not N+1's. This holds for generic targets too,
+ // unlike comparing ISA versions.
+ auto hasGen = [&](unsigned gen) {
+ switch (gen) {
+ case 7:
+ return has(AMDGPU::FEAT_CI_INSTS);
+ case 8:
+ return has(AMDGPU::FEAT_GFX8_INSTS);
+ case 9:
+ return has(AMDGPU::FEAT_GFX9_INSTS);
+ case 10:
+ return has(AMDGPU::FEAT_GFX10_INSTS);
+ case 11:
+ return has(AMDGPU::FEAT_GFX11_INSTS);
+ case 12:
+ return has(AMDGPU::FEAT_GFX12_INSTS);
+ case 13:
+ return has(AMDGPU::FEAT_GFX13_INSTS);
+ default:
+ return false;
+ }
+ };
+
+ if (isUnknown())
+ return false;
+ // gfx6 is the base: it has none of the generation features.
+ if (major == 6)
+ return !hasGen(7);
+ return hasGen(major) && !hasGen(major + 1);
+}
+
+std::optional<unsigned> TargetInfo::getBufferResourceNumRecordsWidth() const {
+ return AMDGPU::getBufferResourceNumRecordsWidth(kind);
+}
+
+std::optional<unsigned> TargetInfo::getMaxAddressableLocalMemorySize() const {
+ if (isUnknown())
+ return std::nullopt;
+ return AMDGPU::getMaxHWAddressableLocalMemorySize(kind);
+}
+
+std::optional<unsigned> TargetInfo::getTotalNumSGPRs() const {
+ if (isUnknown())
+ return std::nullopt;
+ return AMDGPU::getTotalNumSGPRs(kind);
+}
+
+std::optional<unsigned> TargetInfo::getAddressableNumSGPRs() const {
+ if (isUnknown())
+ return std::nullopt;
+ return AMDGPU::getAddressableNumSGPRs(kind);
+}
+
+std::optional<unsigned> TargetInfo::getSGPRAllocGranule() const {
+ if (isUnknown())
+ return std::nullopt;
+ return AMDGPU::getSGPRAllocGranule(kind);
+}
+
+std::optional<unsigned> TargetInfo::getVGPRAllocGranule() const {
+ // The granule depends on the wavefront size, which get() has already pinned.
+ std::optional<unsigned> waveSize = getWavefrontSize();
+ if (isUnknown() || !waveSize)
+ return std::nullopt;
+ return AMDGPU::getVGPRAllocGranule(kind, /*IsWave32=*/*waveSize == 32);
+}
+
+std::optional<unsigned> TargetInfo::getLDSBankCount() const {
+ if (isUnknown())
+ return std::nullopt;
+ return AMDGPU::getLDSBankCount(kind);
+}
+
+std::optional<unsigned> TargetInfo::getMaxWavesPerEU() const {
+ if (isUnknown())
+ return std::nullopt;
+ return AMDGPU::getMaxWavesPerEU(kind);
+}
+
+std::optional<unsigned> TargetInfo::getWavefrontSize() const {
+ if (has(AMDGPU::FEAT_WAVEFRONTSIZE64))
+ return 64;
+ if (has(AMDGPU::FEAT_WAVEFRONTSIZE32))
+ return 32;
+ return std::nullopt;
+}
+
+void TargetInfo::migrateArchFeaturesToModuleFlags(Operation *op) const {
+ assert(LLVM::satisfiesLLVMModule(op) &&
+ "xnack and sramecc describe a whole code object, so they can only be "
+ "recorded on a module");
+ ROCDLDialect *dialect =
+ op->getContext()->getOrLoadDialect<ROCDL::ROCDLDialect>();
+ Builder builder(op->getContext());
+ // The helpers differ in type, hence the generic lambda.
+ auto migrate = [&](AMDGPU::TargetIDSetting setting, auto helper) {
+ if (setting != AMDGPU::TargetIDSetting::On &&
+ setting != AMDGPU::TargetIDSetting::Off)
+ return;
+ helper.setAttr(op,
+ builder.getBoolAttr(setting == AMDGPU::TargetIDSetting::On));
+ };
+ migrate(xnackSetting, dialect->getXnackAttrHelper());
+ migrate(sramEccSetting, dialect->getSrameccAttrHelper());
+}
+
+AMDGPU::IsaVersion TargetInfo::getIsaVersion() const {
+ return AMDGPU::getIsaVersion(subArch);
+}
+
+StringRef TargetInfo::getArchName() const {
+ return AMDGPU::getArchNameAMDGCN(kind);
+}
+
+bool TargetInfo::isGeneric() const {
+ return !isUnknown() && AMDGPU::getMajorSubArch(subArch) == subArch;
+}
diff --git a/mlir/unittests/Dialect/LLVMIR/CMakeLists.txt b/mlir/unittests/Dialect/LLVMIR/CMakeLists.txt
index dc0e82d7591e8..0dffa51bdcf16 100644
--- a/mlir/unittests/Dialect/LLVMIR/CMakeLists.txt
+++ b/mlir/unittests/Dialect/LLVMIR/CMakeLists.txt
@@ -1,9 +1,11 @@
add_mlir_unittest(MLIRLLVMIRTests
LLVMAttrsTest.cpp
LLVMTypeTest.cpp
+ ROCDLTargetInfoTest.cpp
)
mlir_target_link_libraries(MLIRLLVMIRTests
PRIVATE
MLIRFuncDialect
MLIRLLVMDialect
+ MLIRROCDLDialect
)
diff --git a/mlir/unittests/Dialect/LLVMIR/ROCDLTargetInfoTest.cpp b/mlir/unittests/Dialect/LLVMIR/ROCDLTargetInfoTest.cpp
new file mode 100644
index 0000000000000..9528392f5ab86
--- /dev/null
+++ b/mlir/unittests/Dialect/LLVMIR/ROCDLTargetInfoTest.cpp
@@ -0,0 +1,392 @@
+//===- ROCDLTargetInfoTest.cpp - Unit tests for ROCDL::TargetInfo ---------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "mlir/Dialect/LLVMIR/ROCDLTargetInfo.h"
+#include "mlir/Dialect/LLVMIR/ROCDLDialect.h"
+#include "mlir/IR/BuiltinOps.h"
+#include "mlir/IR/Diagnostics.h"
+#include "mlir/IR/MLIRContext.h"
+#include "mlir/IR/OwningOpRef.h"
+#include "gtest/gtest.h"
+
+namespace mlir::ROCDL {
+namespace {
+
+/// Resolves a target, collecting anything reported through emitError.
+FailureOr<TargetInfo> resolve(StringRef arch, unsigned waveSize,
+ std::string &error) {
+ MLIRContext ctx;
+ ScopedDiagnosticHandler handler(&ctx, [&](Diagnostic &diag) {
+ error = diag.str();
+ return success();
+ });
+ return TargetInfo::get(arch, waveSize,
+ [&] { return emitError(UnknownLoc::get(&ctx)); });
+}
+
+/// Resolves a target, asserting that it succeeded.
+TargetInfo getTarget(StringRef arch, unsigned waveSize = 0) {
+ std::string error;
+ FailureOr<TargetInfo> target = resolve(arch, waveSize, error);
+ EXPECT_TRUE(succeeded(target)) << "'" << arch << "': " << error;
+ return succeeded(target) ? *target : TargetInfo();
+}
+
+/// Returns the error message produced when resolving a target, or "" if it
+/// unexpectedly succeeded.
+std::string getTargetError(StringRef arch, unsigned waveSize = 0) {
+ std::string error;
+ FailureOr<TargetInfo> target = resolve(arch, waveSize, error);
+ EXPECT_TRUE(failed(target)) << "expected '" << arch << "' to fail";
+ return error;
+}
+
+TEST(TargetInfoTest, ParseGPUName) {
+ TargetInfo gfx942 = getTarget("gfx942");
+ EXPECT_EQ(gfx942.getArchName(), "gfx942");
+ EXPECT_EQ(gfx942.getSubArch(), llvm::Triple::AMDGPUSubArch942);
+ EXPECT_FALSE(gfx942.isUnknown());
+ EXPECT_FALSE(gfx942.isGeneric());
+
+ llvm::AMDGPU::IsaVersion version = gfx942.getIsaVersion();
+ EXPECT_EQ(version.Major, 9u);
+ EXPECT_EQ(version.Minor, 4u);
+ EXPECT_EQ(version.Stepping, 2u);
+}
+
+TEST(TargetInfoTest, ParseTriple) {
+ // A subarch-bearing triple identifies the GPU on its own.
+ TargetInfo fromTriple = getTarget("amdgpu9.42-amd-amdhsa");
+ EXPECT_EQ(fromTriple.getArchName(), "gfx942");
+ EXPECT_EQ(fromTriple.getSubArch(), llvm::Triple::AMDGPUSubArch942);
+
+ // The legacy subarch-less triple names no GPU.
+ TargetInfo legacy = getTarget("amdgcn-amd-amdhsa");
+ EXPECT_TRUE(legacy.isUnknown());
+ EXPECT_EQ(legacy.getSubArch(), llvm::Triple::NoSubArch);
+ EXPECT_FALSE(legacy.has(llvm::AMDGPU::FEAT_GFX9_INSTS));
+
+ // A full target ID pins the processor down, the way -mcpu does. This is the
+ // spelling rocminfo prints for a device's ISA.
+ TargetInfo full = getTarget("amdgcn-amd-amdhsa--gfx942");
+ EXPECT_EQ(full.getArchName(), "gfx942");
+ EXPECT_EQ(full.getSubArch(), llvm::Triple::AMDGPUSubArch942);
+
+ // A named environment parses too.
+ EXPECT_EQ(getTarget("amdgcn-amd-amdhsa-unknown-gfx942").getArchName(),
+ "gfx942");
+
+ // A family triple plus a processor narrows to the exact GPU.
+ TargetInfo family = getTarget("amdgpu9.4-amd-amdhsa--gfx950");
+ EXPECT_EQ(family.getArchName(), "gfx950");
+ EXPECT_EQ(family.getSubArch(), llvm::Triple::AMDGPUSubArch950);
+}
+
+TEST(TargetInfoTest, ParseGeneric) {
+ // Generic targets are representable, and carry only the features common to
+ // every GPU they cover.
+ TargetInfo generic = getTarget("gfx9-4-generic");
+ EXPECT_TRUE(generic.isGeneric());
+ EXPECT_EQ(generic.getArchName(), "gfx9-4-generic");
+
+ // gfx9-4-generic covers gfx942 and gfx950, so it must not claim anything
+ // exclusive to either.
+ EXPECT_TRUE(generic.has(llvm::AMDGPU::FEAT_GFX940_INSTS));
+ EXPECT_FALSE(generic.has(llvm::AMDGPU::FEAT_GFX950_INSTS));
+ EXPECT_FALSE(generic.has(llvm::AMDGPU::FEAT_XF32_INSTS));
+ EXPECT_TRUE(getTarget("gfx942").has(llvm::AMDGPU::FEAT_XF32_INSTS));
+}
+
+TEST(TargetInfoTest, ParseInvalid) {
+ // Unlike Chipset::parse, a well-formed but nonexistent GPU is rejected.
+ EXPECT_NE(getTargetError("gfx999"), "");
+ EXPECT_NE(getTargetError("gfx000"), "");
+ EXPECT_NE(getTargetError("navi33"), "");
+ EXPECT_NE(getTargetError("sm_80"), "");
+ EXPECT_NE(getTargetError("GFX942"), "");
+ EXPECT_NE(getTargetError(""), "");
+
+ // Triple parsing maps any unrecognized "amdgpu..." to NoSubArch rather than
+ // reporting an error, so a typo must not be mistaken for the legacy triple.
+ EXPECT_NE(getTargetError("amdgpu9.99-amd-amdhsa"), "");
+ EXPECT_NE(getTargetError("amdgputypo-amd-amdhsa"), "");
+
+ // A processor inconsistent with the triple's subarch is rejected.
+ EXPECT_NE(getTargetError("amdgpu9.42-amd-amdhsa--gfx1030"), "");
+ EXPECT_NE(getTargetError("amdgcn-amd-amdhsa--gfx999"), "");
+}
+
+TEST(TargetInfoTest, TargetIDModifiers) {
+ using llvm::AMDGPU::TargetIDSetting;
+
+ // xnack and sramecc are the only modifiers a target ID admits, matching
+ // clang::parseTargetID. Unspecified means "either".
+ EXPECT_EQ(getTarget("gfx90a").getXnackSetting(), TargetIDSetting::Any);
+ EXPECT_EQ(getTarget("gfx90a:xnack+").getXnackSetting(), TargetIDSetting::On);
+ EXPECT_EQ(getTarget("gfx90a:xnack-").getXnackSetting(), TargetIDSetting::Off);
+ EXPECT_EQ(getTarget("gfx90a:sramecc+").getSramEccSetting(),
+ TargetIDSetting::On);
+ EXPECT_EQ(getTarget("gfx90a:sramecc-:xnack+").getSramEccSetting(),
+ TargetIDSetting::Off);
+
+ // A GPU that cannot switch them reports Unsupported and rejects a modifier.
+ EXPECT_EQ(getTarget("gfx600").getXnackSetting(),
+ TargetIDSetting::Unsupported);
+ EXPECT_NE(getTargetError("gfx600:xnack+"), "");
+
+ // The sign is mandatory, the feature must be known, and nothing else is a
+ // target-ID feature -- in particular wavefront size cannot ride on `arch`.
+ EXPECT_NE(getTargetError("gfx908:xnack"), "");
+ EXPECT_NE(getTargetError("gfx942:not-a-feature+"), "");
+ EXPECT_NE(getTargetError("gfx1030:wavefrontsize64+"), "");
+
+ // Modifiers do not disturb the feature set.
+ EXPECT_TRUE(getTarget("gfx90a:xnack+").has(llvm::AMDGPU::FEAT_MAI_INSTS));
+}
+
+TEST(TargetInfoTest, WavefrontSize) {
+ // Single-mode targets report their only size.
+ EXPECT_EQ(getTarget("gfx90a").getWavefrontSize(), 64u);
+ EXPECT_EQ(getTarget("gfx942").getWavefrontSize(), 64u);
+ EXPECT_EQ(getTarget("gfx1250").getWavefrontSize(), 32u);
+
+ // Targets supporting both default to wave32, and honour an explicit request.
+ // This is the case a triple alone cannot express.
+ for (StringRef gpu : {"gfx1030", "gfx1100", "gfx1200"}) {
+ EXPECT_EQ(getTarget(gpu).getWavefrontSize(), 32u) << gpu;
+ EXPECT_EQ(getTarget(gpu, /*waveSize=*/64).getWavefrontSize(), 64u) << gpu;
+ EXPECT_EQ(getTarget(gpu, /*waveSize=*/32).getWavefrontSize(), 32u) << gpu;
+ }
+
+ // Naming the size a single-mode target already runs at is accepted.
+ EXPECT_EQ(getTarget("gfx942", /*waveSize=*/64).getWavefrontSize(), 64u);
+ EXPECT_EQ(getTarget("gfx1250", /*waveSize=*/32).getWavefrontSize(), 32u);
+
+ // Asking it for the other size is an error, not a silent mis-lowering.
+ EXPECT_NE(getTargetError("gfx942", /*waveSize=*/32), "");
+ EXPECT_NE(getTargetError("gfx1250", /*waveSize=*/64), "");
+
+ // Only 32 and 64 are wavefront sizes.
+ EXPECT_NE(getTargetError("gfx1030", /*waveSize=*/17), "");
+ EXPECT_NE(getTargetError("gfx1030", /*waveSize=*/128), "");
+
+ // An unknown target has no wavefront size, but still rejects nonsense.
+ EXPECT_EQ(getTarget("amdgcn-amd-amdhsa").getWavefrontSize(), std::nullopt);
+ EXPECT_NE(getTargetError("amdgcn-amd-amdhsa", /*waveSize=*/17), "");
+}
+
+TEST(TargetInfoTest, SupportsBothWavefrontSizes) {
+ // Only these leave the choice to the features; the rest pin a size, which is
+ // why the pipeline's wave64 option must not be forced onto them.
+ for (StringRef gpu : {"gfx1030", "gfx1100", "gfx1200"})
+ EXPECT_TRUE(getTarget(gpu).supportsBothWavefrontSizes()) << gpu;
+ for (StringRef gpu : {"gfx90a", "gfx942", "gfx1250"})
+ EXPECT_FALSE(getTarget(gpu).supportsBothWavefrontSizes()) << gpu;
+
+ // Naming a size does not change what the GPU is capable of.
+ EXPECT_TRUE(
+ getTarget("gfx1030", /*waveSize=*/64).supportsBothWavefrontSizes());
+ EXPECT_FALSE(getTarget("amdgcn-amd-amdhsa").supportsBothWavefrontSizes());
+}
+
+TEST(TargetInfoTest, Fp8Formats) {
+ // hasFnuzFp8 is not the negation of hasOcpFp8: a target with no fp8
+ // conversions at all has neither.
+ EXPECT_TRUE(getTarget("gfx942").hasFnuzFp8());
+ EXPECT_FALSE(getTarget("gfx942").hasOcpFp8());
+ for (StringRef gpu : {"gfx950", "gfx1170", "gfx1200"}) {
+ EXPECT_FALSE(getTarget(gpu).hasFnuzFp8()) << gpu;
+ EXPECT_TRUE(getTarget(gpu).hasOcpFp8()) << gpu;
+ }
+ for (StringRef gpu : {"gfx908", "gfx90a", "gfx900"}) {
+ EXPECT_FALSE(getTarget(gpu).hasFnuzFp8()) << gpu;
+ EXPECT_FALSE(getTarget(gpu).hasOcpFp8()) << gpu;
+ }
+}
+
+TEST(TargetInfoTest, BufferResourceNumRecordsWidth) {
+ // A width rather than a capability bit, so that a third width would not need
+ // every caller to learn a new predicate.
+ for (StringRef gpu : {"gfx900", "gfx1030", "gfx1200", "gfx1201"})
+ EXPECT_EQ(getTarget(gpu).getBufferResourceNumRecordsWidth(), 32u) << gpu;
+ for (StringRef gpu : {"gfx1250", "gfx1251", "gfx1250-strict"})
+ EXPECT_EQ(getTarget(gpu).getBufferResourceNumRecordsWidth(), 45u) << gpu;
+
+ // Generic targets take the width of the family they cover.
+ EXPECT_EQ(getTarget("gfx12-generic").getBufferResourceNumRecordsWidth(), 32u);
+ EXPECT_EQ(getTarget("gfx12-5-generic").getBufferResourceNumRecordsWidth(),
+ 45u);
+
+ // An unknown target has no width, so a lowering that needs one must bail
+ // rather than assume the narrow case.
+ EXPECT_EQ(getTarget("amdgcn-amd-amdhsa").getBufferResourceNumRecordsWidth(),
+ std::nullopt);
+ EXPECT_EQ(TargetInfo().getBufferResourceNumRecordsWidth(), std::nullopt);
+}
+
+TEST(TargetInfoTest, MaxAddressableLocalMemorySize) {
+ EXPECT_EQ(getTarget("gfx900").getMaxAddressableLocalMemorySize(), 65536u);
+ EXPECT_EQ(getTarget("gfx1030").getMaxAddressableLocalMemorySize(), 65536u);
+ EXPECT_EQ(getTarget("gfx950").getMaxAddressableLocalMemorySize(), 163840u);
+ EXPECT_EQ(getTarget("gfx1250").getMaxAddressableLocalMemorySize(), 327680u);
+
+ EXPECT_EQ(getTarget("amdgcn-amd-amdhsa").getMaxAddressableLocalMemorySize(),
+ std::nullopt);
+}
+
+TEST(TargetInfoTest, RegisterAndLDSProperties) {
+ // These forward to the TargetParser, so rather than restating its tables,
+ // check that each agrees with the function it wraps and that an unresolved
+ // target reports "unknown" instead of the LLVM-side hardcoded fallback.
+ for (StringRef gpu : {"gfx900", "gfx90a", "gfx942", "gfx1030", "gfx1250"}) {
+ TargetInfo target = getTarget(gpu);
+ llvm::AMDGPU::GPUKind kind = target.getGPUKind();
+ EXPECT_EQ(target.getTotalNumSGPRs(), llvm::AMDGPU::getTotalNumSGPRs(kind))
+ << gpu;
+ EXPECT_EQ(target.getAddressableNumSGPRs(),
+ llvm::AMDGPU::getAddressableNumSGPRs(kind))
+ << gpu;
+ EXPECT_EQ(target.getSGPRAllocGranule(),
+ llvm::AMDGPU::getSGPRAllocGranule(kind))
+ << gpu;
+ EXPECT_EQ(target.getLDSBankCount(), llvm::AMDGPU::getLDSBankCount(kind))
+ << gpu;
+ EXPECT_EQ(target.getMaxWavesPerEU(), llvm::AMDGPU::getMaxWavesPerEU(kind))
+ << gpu;
+ }
+
+ // The VGPR granule depends on the wavefront size, which TargetInfo supplies
+ // from the resolved target rather than taking as an argument.
+ TargetInfo wave32 = getTarget("gfx1030", /*waveSize=*/32);
+ TargetInfo wave64 = getTarget("gfx1030", /*waveSize=*/64);
+ EXPECT_EQ(wave32.getVGPRAllocGranule(),
+ llvm::AMDGPU::getVGPRAllocGranule(wave32.getGPUKind(),
+ /*IsWave32=*/true));
+ EXPECT_EQ(wave64.getVGPRAllocGranule(),
+ llvm::AMDGPU::getVGPRAllocGranule(wave64.getGPUKind(),
+ /*IsWave32=*/false));
+
+ TargetInfo unknown = getTarget("amdgcn-amd-amdhsa");
+ EXPECT_EQ(unknown.getTotalNumSGPRs(), std::nullopt);
+ EXPECT_EQ(unknown.getAddressableNumSGPRs(), std::nullopt);
+ EXPECT_EQ(unknown.getSGPRAllocGranule(), std::nullopt);
+ EXPECT_EQ(unknown.getVGPRAllocGranule(), std::nullopt);
+ EXPECT_EQ(unknown.getLDSBankCount(), std::nullopt);
+ EXPECT_EQ(unknown.getMaxWavesPerEU(), std::nullopt);
+ EXPECT_EQ(TargetInfo().getLDSBankCount(), std::nullopt);
+}
+
+TEST(TargetInfoTest, Generation) {
+ EXPECT_TRUE(getTarget("gfx900").isGeneration(9));
+ EXPECT_TRUE(getTarget("gfx942").isGeneration(9));
+ EXPECT_TRUE(getTarget("gfx950").isGeneration(9));
+ EXPECT_FALSE(getTarget("gfx942").isGeneration(10));
+ EXPECT_FALSE(getTarget("gfx942").isGeneration(8));
+
+ EXPECT_TRUE(getTarget("gfx1010").isGeneration(10));
+ EXPECT_TRUE(getTarget("gfx1030").isGeneration(10));
+ EXPECT_TRUE(getTarget("gfx1100").isGeneration(11));
+ EXPECT_TRUE(getTarget("gfx1200").isGeneration(12));
+ EXPECT_TRUE(getTarget("gfx1250").isGeneration(12));
+ EXPECT_TRUE(getTarget("gfx803").isGeneration(8));
+ EXPECT_TRUE(getTarget("gfx700").isGeneration(7));
+ EXPECT_TRUE(getTarget("gfx600").isGeneration(6));
+
+ // Generic targets report the generation of the family they cover, which
+ // comparing ISA versions would get wrong.
+ EXPECT_TRUE(getTarget("gfx9-4-generic").isGeneration(9));
+ EXPECT_TRUE(getTarget("gfx11-generic").isGeneration(11));
+ EXPECT_TRUE(getTarget("gfx12-generic").isGeneration(12));
+
+ // An unknown target is in no generation.
+ EXPECT_FALSE(getTarget("amdgcn-amd-amdhsa").isGeneration(9));
+}
+
+/// The two module attributes `migrateArchFeaturesToModuleFlags` writes, as read
+/// back off the module; a null attribute means it wrote nothing.
+struct MigratedSettings {
+ BoolAttr xnack;
+ BoolAttr sramecc;
+};
+
+/// Migrates `arch`'s target-ID modifiers onto a fresh module, after seeding it
+/// with whatever `preset` puts there, and reads the result back.
+MigratedSettings
+migrate(StringRef arch,
+ function_ref<void(ROCDLDialect *, Operation *)> preset = nullptr) {
+ MLIRContext ctx;
+ auto *dialect = ctx.getOrLoadDialect<ROCDLDialect>();
+ OwningOpRef<ModuleOp> module = ModuleOp::create(UnknownLoc::get(&ctx));
+ if (preset)
+ preset(dialect, *module);
+ getTarget(arch).migrateArchFeaturesToModuleFlags(*module);
+ return {dialect->getXnackAttrHelper().getAttr(*module),
+ dialect->getSrameccAttrHelper().getAttr(*module)};
+}
+
+TEST(TargetInfoTest, MigrateArchFeaturesToModuleFlags) {
+ // A pinned modifier becomes the matching attribute, either way round.
+ EXPECT_TRUE(migrate("gfx90a:xnack+").xnack.getValue());
+ EXPECT_FALSE(migrate("gfx90a:xnack-").xnack.getValue());
+ EXPECT_FALSE(migrate("gfx90a:sramecc-").sramecc.getValue());
+
+ MigratedSettings both = migrate("gfx90a:sramecc+:xnack-");
+ EXPECT_TRUE(both.sramecc.getValue());
+ EXPECT_FALSE(both.xnack.getValue());
+
+ // Only the pinned one is written: the other is still "either".
+ EXPECT_EQ(migrate("gfx90a:xnack+").sramecc, nullptr);
+
+ // A target ID that pins neither, and a GPU that has neither to pin, write
+ // nothing at all rather than writing false.
+ MigratedSettings any = migrate("gfx90a");
+ EXPECT_EQ(any.xnack, nullptr);
+ EXPECT_EQ(any.sramecc, nullptr);
+ EXPECT_EQ(migrate("gfx600").xnack, nullptr);
+
+ // An `arch` that says nothing leaves an explicit setting alone, while one
+ // that does pin the setting is the more specific request and overrides it.
+ auto presetXnackTrue = [](ROCDLDialect *dialect, Operation *op) {
+ dialect->getXnackAttrHelper().setAttr(
+ op, BoolAttr::get(op->getContext(), true));
+ };
+ EXPECT_TRUE(migrate("gfx90a", presetXnackTrue).xnack.getValue());
+ EXPECT_FALSE(migrate("gfx90a:xnack-", presetXnackTrue).xnack.getValue());
+}
+
+TEST(TargetInfoTest, ResolveArchOption) {
+ // Anything actually passed as `arch` wins over the deprecated alias, whether
+ // or not the alias was given.
+ EXPECT_EQ(resolveArchOption("gfx942", ""), "gfx942");
+ EXPECT_EQ(resolveArchOption("gfx942", "gfx90a"), "gfx942");
+
+ // The alias is consulted only for the sentinels that mean "no target given":
+ // "invalid" for the passes that require one, and "" for the passes where
+ // asking for no target is a legitimate request.
+ EXPECT_EQ(resolveArchOption("invalid", "gfx90a"), "gfx90a");
+ EXPECT_EQ(resolveArchOption("", "gfx90a"), "gfx90a");
+
+ // A target ID goes through the alias unchanged, so an old invocation that
+ // named one keeps naming it.
+ EXPECT_EQ(resolveArchOption("invalid", "gfx90a:xnack+"), "gfx90a:xnack+");
+
+ // With neither given, the sentinel survives, so the pass still reports it as
+ // the target it could not resolve rather than reporting an empty name.
+ EXPECT_EQ(resolveArchOption("invalid", ""), "invalid");
+ EXPECT_EQ(resolveArchOption("", ""), "");
+}
+
+TEST(TargetInfoTest, DefaultIsUnknown) {
+ TargetInfo target;
+ EXPECT_TRUE(target.isUnknown());
+ EXPECT_FALSE(target.has(llvm::AMDGPU::FEAT_GFX9_INSTS));
+ EXPECT_EQ(target.getArchName(), "");
+ EXPECT_EQ(target.getWavefrontSize(), std::nullopt);
+}
+} // namespace
+} // namespace mlir::ROCDL
>From ba2e1dfcd66c8c05e8ba0826881c3a54bb3adeea Mon Sep 17 00:00:00 2001
From: Krzysztof Drewniak <Krzysztof.Drewniak at amd.com>
Date: Thu, 17 Sep 2026 16:02:54 +0000
Subject: [PATCH 2/2] Drop namespacing
---
mlir/unittests/Dialect/LLVMIR/ROCDLTargetInfoTest.cpp | 5 +++--
1 file changed, 3 insertions(+), 2 deletions(-)
diff --git a/mlir/unittests/Dialect/LLVMIR/ROCDLTargetInfoTest.cpp b/mlir/unittests/Dialect/LLVMIR/ROCDLTargetInfoTest.cpp
index 9528392f5ab86..2818418ddecc1 100644
--- a/mlir/unittests/Dialect/LLVMIR/ROCDLTargetInfoTest.cpp
+++ b/mlir/unittests/Dialect/LLVMIR/ROCDLTargetInfoTest.cpp
@@ -14,7 +14,9 @@
#include "mlir/IR/OwningOpRef.h"
#include "gtest/gtest.h"
-namespace mlir::ROCDL {
+using namespace mlir;
+using namespace mlir::ROCDL;
+
namespace {
/// Resolves a target, collecting anything reported through emitError.
@@ -389,4 +391,3 @@ TEST(TargetInfoTest, DefaultIsUnknown) {
EXPECT_EQ(target.getWavefrontSize(), std::nullopt);
}
} // namespace
-} // namespace mlir::ROCDL
More information about the Mlir-commits
mailing list