[llvm-branch-commits] [mlir] [mlir][ROCDL] Add TargetInfo to replace Chipset, allow features queries (PR #223562)
via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Mon Sep 14 16:49:31 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-mlir
Author: Krzysztof Drewniak (krzysz00)
<details>
<summary>Changes</summary>
Add a new 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.
---
Patch is 35.59 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/223562.diff
5 Files Affected:
- (added) mlir/include/mlir/Dialect/LLVMIR/ROCDLTargetInfo.h (+193)
- (modified) mlir/lib/Dialect/LLVMIR/CMakeLists.txt (+2)
- (added) mlir/lib/Dialect/LLVMIR/IR/ROCDLTargetInfo.cpp (+242)
- (modified) mlir/unittests/Dialect/LLVMIR/CMakeLists.txt (+2)
- (added) mlir/unittests/Dialect/LLVMIR/ROCDLTargetInfoTest.cpp (+392)
``````````diff
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...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/223562
More information about the llvm-branch-commits
mailing list