[llvm-branch-commits] [mlir] [mlir] Migrate AMDGPU/ROCDL to targets, not chipset versions (PR #223563)
Krzysztof Drewniak via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Thu Sep 17 09:03:51 PDT 2026
https://github.com/krzysz00 updated https://github.com/llvm/llvm-project/pull/223563
>From 04d87a2d29d2d0b2f08fcefde1ccba1fabe004a4 Mon Sep 17 00:00:00 2001
From: Krzysztof Drewniak <Krzysztof.Drewniak at amd.com>
Date: Mon, 31 Aug 2026 19:32:37 +0000
Subject: [PATCH] [mlir] Migrate AMDGPU/ROCDL to targets, not chipset versions
**migration tl;dr:** Replace usages of `amdgpu::Chipset` with `ROCDL::TargetInfo`, ideally move from `chipset=` to `arch=`. If you don't use upstream pipelines, call 'TargetInfo::migrateArchFeaturesToModuleFlags` at the appropriate location.
Further note: if you've got a build pipeline that's getting a `gfxXXX` name from something like `rocm_agent_enumerator`, using a full triple name like the ones you get from `rocminfo` is preferred.
`amdgpu::Chipset` was an awkward hack that was hard to keep up to date
with changes in the compiler/new architectures, and didn't properly
support generic targets (and has been strongly disfavored by the
compiler team).
This PR replaces `amdgpu::Chipset` with `ROCDL::TargetInfo`, a
structure that uses LLVM's TargetParser and the underlying LLVM
features tables to get the real nature of the target being compiled
for.
This also helps MLIR move to
new-style (`-mtriple=amdgpuX.YZ-amd-amdhsa`) over "old
style" (`-mtriple=amdgcn-amd-amdhsa -mcpu=gfxXYZ`) triples.
The utility structure is moved from AMDGPU to ROCDL, both because it's
tied to LLVM rather directly and because projects like Triton should
be able to use these feature tests without pulling in the AMDGPU
dialect and its memref dependencies.
This migration also fixes a few correctness issues:
- Atomic emulation was producing floating-point additions that don't
exist on gfx90c (even though it's "after" gfx90a) and gfx908's more
precise about where emulation is needed.
- gfx90c was also being handed a bare `s_barrier`, but it has no
hardware barrier back-off, so it needs the inline asm workaround.
- gfx950 won't allow xf32 MFMAs anymore.
- permlane_swap forms that don't exist on some architectures no
longer lower.
- gfx11.7 is now listed as an OCP FP8-having target.
Some checks still need to check for a generation (ex. when encoding
s_waitcnt or what the semantics of a WMMA are) go through an
`isGeneration(N)` method, which checks for having instructions from
generation N but not N+1.
(The barrier lowering has been reordered to account for gfx13 not
having, but also not needing, BackOffBarrier.)
This migration renames the `chipset` or `chip` options on most passes to `arch`, but keeps the old name for compatibility.
Please note that, because of changes to how LLVM handles xnack and sramecc (they aren't target features anymore), you need to set `rocdl.xnack` and `rocdl.sramecc` on the modules you're translating if you know the values of those modifiers. `TargetInfo::migrateArchFeaturesToModuleFlags(Operation *op)` will do this for you.
Chipset is deprecated instead of being removed so that folks have time
to migrate.
Pre-commit tests in #220104.
AI disclosure: I steered this, Claude wrote the code, I tried to clean
up the docs.
---
mlir/docs/ReleaseNotes.md | 33 +
.../Conversion/AMDGPUToROCDL/AMDGPUToROCDL.h | 4 +-
.../Conversion/ArithToAMDGPU/ArithToAMDGPU.h | 4 +-
.../Conversion/GPUToROCDL/GPUToROCDLPass.h | 7 +-
.../mlir/Conversion/MathToROCDL/MathToROCDL.h | 8 +-
mlir/include/mlir/Conversion/Passes.td | 46 +-
.../mlir/Dialect/AMDGPU/Transforms/Passes.h | 4 +-
.../mlir/Dialect/AMDGPU/Transforms/Passes.td | 16 +-
.../mlir/Dialect/AMDGPU/Utils/Chipset.h | 12 +
.../mlir/Dialect/GPU/Pipelines/Passes.h | 29 +-
.../GPU/TransformOps/GPUTransformOps.td | 27 +-
.../mlir/Dialect/GPU/Transforms/Passes.h | 15 +-
.../mlir/Dialect/GPU/Transforms/Passes.td | 19 +-
.../AMDGPUToROCDL/AMDGPUToROCDL.cpp | 639 +++++++++---------
.../ArithToAMDGPU/ArithToAMDGPU.cpp | 48 +-
.../GPUToROCDL/LowerGpuOpsToROCDLOps.cpp | 60 +-
.../Conversion/MathToROCDL/MathToROCDL.cpp | 27 +-
.../Dialect/AMDGPU/Transforms/CMakeLists.txt | 1 +
.../AMDGPU/Transforms/EmulateAtomics.cpp | 86 +--
.../GPU/Pipelines/GPUToROCDLPipeline.cpp | 14 +-
.../GPU/TransformOps/GPUTransformOps.cpp | 45 +-
.../GPU/Transforms/PromoteShuffleToAMDGPU.cpp | 8 +-
.../GPU/Transforms/ROCDLAttachTarget.cpp | 91 ++-
.../GPU/Transforms/SubgroupReduceLowering.cpp | 34 +-
.../8-bit-floats-ocp-gfx1170.mlir | 23 -
.../AMDGPUToROCDL/8-bit-floats-ocp.mlir | 5 +-
.../AMDGPUToROCDL/8-bit-floats.mlir | 2 +-
.../AMDGPUToROCDL/amdgpu-to-rocdl.mlir | 14 +-
.../Conversion/AMDGPUToROCDL/dot-gfx11.mlir | 2 +-
.../Conversion/AMDGPUToROCDL/dot-gfx12.mlir | 2 +-
.../Conversion/AMDGPUToROCDL/dot-gfx9.mlir | 2 +-
.../Conversion/AMDGPUToROCDL/dot-invalid.mlir | 4 +-
mlir/test/Conversion/AMDGPUToROCDL/dpp.mlir | 7 +-
.../Conversion/AMDGPUToROCDL/gfx1250.mlir | 2 +-
.../AMDGPUToROCDL/global-prefetch.mlir | 2 +-
.../AMDGPUToROCDL/global_transpose_load.mlir | 6 +-
.../AMDGPUToROCDL/lds-barrier-gfx90c.mlir | 5 +-
.../AMDGPUToROCDL/load_lds-gfx950.mlir | 4 +-
.../Conversion/AMDGPUToROCDL/load_lds.mlir | 4 +-
.../AMDGPUToROCDL/memory_counter_wait.mlir | 8 +-
.../memory_counter_wait_tensor.mlir | 2 +-
.../memory_counter_wait_unsupported.mlir | 6 +-
.../AMDGPUToROCDL/mfma-fp8-invalid.mlir | 21 +
.../Conversion/AMDGPUToROCDL/mfma-gfx950.mlir | 13 +-
.../mfma-reduce-precision-invalid.mlir | 23 +
mlir/test/Conversion/AMDGPUToROCDL/mfma.mlir | 2 +-
.../Conversion/AMDGPUToROCDL/packed-ext.mlir | 2 +-
.../AMDGPUToROCDL/packed-trunc-invalid.mlir | 2 +-
.../AMDGPUToROCDL/packed-trunc.mlir | 2 +-
.../permlane-gfx1200-invalid.mlir | 21 +
.../permlane-gfx1250-invalid.mlir | 11 +
.../AMDGPUToROCDL/permlane-gfx1250.mlir | 12 +
.../AMDGPUToROCDL/permlane-var.mlir | 2 +-
.../Conversion/AMDGPUToROCDL/permlane.mlir | 3 +-
.../AMDGPUToROCDL/sparse-mfma-gfx950.mlir | 2 +-
.../Conversion/AMDGPUToROCDL/sparse-mfma.mlir | 2 +-
.../Conversion/AMDGPUToROCDL/swizzle.mlir | 2 +-
.../AMDGPUToROCDL/swmmac-gfx12.mlir | 2 +-
.../AMDGPUToROCDL/swmmac-gfx1250.mlir | 2 +-
.../AMDGPUToROCDL/transpose_load.mlir | 4 +-
.../AMDGPUToROCDL/transpose_load_gfx1250.mlir | 2 +-
.../transpose_load_gfx1250_invalid.mlir | 2 +-
.../transpose_load_gfx950_invalid.mlir | 2 +-
.../AMDGPUToROCDL/transpose_load_reject.mlir | 2 +-
.../Conversion/AMDGPUToROCDL/wmma-gfx11.mlir | 2 +-
.../Conversion/AMDGPUToROCDL/wmma-gfx12.mlir | 2 +-
.../AMDGPUToROCDL/wmma-gfx1250.mlir | 2 +-
.../ArithToAMDGPU/16-bit-floats.mlir | 2 +-
.../8-bit-float-saturation-ocp.mlir | 4 +-
.../ArithToAMDGPU/8-bit-float-saturation.mlir | 2 +-
.../ArithToAMDGPU/8-bit-floats-ocp.mlir | 6 +-
.../ArithToAMDGPU/8-bit-floats.mlir | 2 +-
.../deprecated-chipset-alias.mlir | 27 +
.../ArithToAMDGPU/scaling-extf.mlir | 4 +-
.../ArithToAMDGPU/scaling-truncf-tensor.mlir | 2 +-
.../ArithToAMDGPU/scaling-truncf.mlir | 4 +-
.../Conversion/GPUCommon/lower-global-id.mlir | 2 +-
.../GPUCommon/lower-memory-space-attrs.mlir | 2 +-
.../GPUCommon/memory-attrbution.mlir | 2 +-
.../GPUCommon/memref-arg-attrs.mlir | 2 +-
.../GPUCommon/memref-arg-noalias-attrs.mlir | 2 +-
.../GPUCommon/memref-arg-noalias-warning.mlir | 2 +-
.../GPUToROCDL/constant-address-space.mlir | 2 +-
.../GPUToROCDL/gpu-to-rocdl-barrier.mlir | 4 +-
.../gpu-to-rocdl-barriers-gfx12.mlir | 2 +-
.../GPUToROCDL/gpu-to-rocdl-hip.mlir | 4 +-
.../gpu-to-rocdl-invalid-ballot.mlir | 2 +-
.../gpu-to-rocdl-invalid-dialect.mlir | 2 +-
.../gpu-to-rocdl-invalid-named-barrier.mlir | 2 +-
.../gpu-to-rocdl-named-barrier-non-const.mlir | 2 +-
.../GPUToROCDL/gpu-to-rocdl-opencl.mlir | 2 +-
.../GPUToROCDL/gpu-to-rocdl-subgroup-id.mlir | 4 +-
.../Conversion/GPUToROCDL/gpu-to-rocdl.mlir | 6 +-
mlir/test/Conversion/GPUToROCDL/memref.mlir | 4 +-
.../private-func-no-c-interface.mlir | 2 +-
.../Conversion/MathToROCDL/math-to-rocdl.mlir | 4 +-
.../AMDGPU/amdgpu-emulate-atomics.mlir | 73 +-
.../GPU/promote-shuffle-amdgpu-invalid.mlir | 58 ++
.../Dialect/GPU/promote-shuffle-amdgpu.mlir | 2 +-
mlir/test/Dialect/LLVMIR/attach-targets.mlir | 2 +-
.../LLVMIR/rocdl-attach-target-arch.mlir | 105 +++
.../GPU/ROCM/gpu-lower-to-rocdl-pipeline.mlir | 2 +-
.../Integration/GPU/ROCM/gpu-to-hsaco.mlir | 2 +-
mlir/test/Integration/GPU/ROCM/printf.mlir | 2 +-
.../Integration/GPU/ROCM/two-modules.mlir | 2 +-
mlir/test/Integration/GPU/ROCM/vecadd.mlir | 2 +-
.../GPU/ROCM/vector-transferops.mlir | 2 +-
mlir/test/lib/Dialect/GPU/TestGpuRewrite.cpp | 13 +-
.../Dialect/AMDGPU/AMDGPUUtilsTest.cpp | 8 +
109 files changed, 1181 insertions(+), 718 deletions(-)
delete mode 100644 mlir/test/Conversion/AMDGPUToROCDL/8-bit-floats-ocp-gfx1170.mlir
create mode 100644 mlir/test/Conversion/AMDGPUToROCDL/mfma-fp8-invalid.mlir
create mode 100644 mlir/test/Conversion/AMDGPUToROCDL/mfma-reduce-precision-invalid.mlir
create mode 100644 mlir/test/Conversion/AMDGPUToROCDL/permlane-gfx1200-invalid.mlir
create mode 100644 mlir/test/Conversion/AMDGPUToROCDL/permlane-gfx1250-invalid.mlir
create mode 100644 mlir/test/Conversion/AMDGPUToROCDL/permlane-gfx1250.mlir
create mode 100644 mlir/test/Conversion/ArithToAMDGPU/deprecated-chipset-alias.mlir
create mode 100644 mlir/test/Dialect/GPU/promote-shuffle-amdgpu-invalid.mlir
create mode 100644 mlir/test/Dialect/LLVMIR/rocdl-attach-target-arch.mlir
diff --git a/mlir/docs/ReleaseNotes.md b/mlir/docs/ReleaseNotes.md
index 16b93c8909670..b4650b82841b3 100644
--- a/mlir/docs/ReleaseNotes.md
+++ b/mlir/docs/ReleaseNotes.md
@@ -8,6 +8,39 @@ specifically, it is a snapshot of the MLIR development at the time of the releas
[TOC]
+## LLVM 24
+
+### GPU/AMDGPU Changes
+
+- `mlir::amdgpu::Chipset` is deprecated in favour of `mlir::ROCDL::TargetInfo`,
+ which describes a target by its triple, subarch, and the resolved set of
+ target features from LLVM's own tables. Lowerings should ask whether a target
+ has a feature rather than inaccurately compare chipset versions.
+ `TargetInfo` also represents generic targets such as `gfx9-4-generic` and, unlike
+ `Chipset`, explicitly stores the wavesize for targets where it is configurable.
+- The `chipset` option in AMDGPU passes is renamed to an `arch` option, which uses
+ Clang target naming syntax. It accepts a GPU name with optional
+ modifiers (`gfx942`, `gfx942:xnack+`, `gfx9-4-generic`), a triple
+ (`amdgpu9.42-amd-amdhsa`), or a full target ID
+ (`amdgpu9.42-amd-amdhsa--gfx90a:sramecc+:xnack-`, which is what `rocminfo` prints
+ for a device's ISA). `chipset` or `chip` remain as compatibility names.
+ The default arch is `invalid`, so a target must be passed
+ explicitly, removing the old "fallback" `gfx000` GPU.
+- Wavefront size is not a target-ID feature, so `convert-gpu-to-rocdl` takes it
+ as a separate `wavesize` option (32, 64, or 0 for the architecture's
+ default). The `wave64` flag on `gpu-lower-to-rocdl-pipeline` and on
+ `rocdl-attach-target` is likewise replaced by the same `wavesize` option,
+ which has the same allowed values.
+- In keeping with broader LLVM changes, `xnack` and `sramecc` are no longer
+ architecture features but module flags. In keeping with Clang, the target
+ specifier still includes these xnack/sramecc flags where they're configurable,
+ but lowering passes now convert these to module flags. Downstream users should
+ call `migrateArchFeaturesToModuleFlags` to lower these attributes in custom
+ pipelines.
+- `rocdl-attach-target` gains `arch` alongside its existing `triple`, `chip` and
+ `features`. When `arch` is given, it overrides `triple` and `chip`, and handles
+ xnack/sramecc modifier migration.
+
## LLVM 21
### GPU/NVVM Changes
diff --git a/mlir/include/mlir/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.h b/mlir/include/mlir/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.h
index 393658652dbac..861bbe0cc8f75 100644
--- a/mlir/include/mlir/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.h
+++ b/mlir/include/mlir/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.h
@@ -8,7 +8,7 @@
#ifndef MLIR_CONVERSION_AMDGPUTOROCDL_AMDGPUTOROCDL_H_
#define MLIR_CONVERSION_AMDGPUTOROCDL_AMDGPUTOROCDL_H_
-#include "mlir/Dialect/AMDGPU/Utils/Chipset.h"
+#include "mlir/Dialect/LLVMIR/ROCDLTargetInfo.h"
#include <memory>
#include <string>
@@ -27,7 +27,7 @@ class Pass;
/// populateAMDGPUTypeAndAttributeConversions().
void populateAMDGPUToROCDLConversionPatterns(LLVMTypeConverter &converter,
RewritePatternSet &patterns,
- amdgpu::Chipset chipset);
+ const ROCDL::TargetInfo &target);
namespace amdgpu {
/// Remap common GPU memory spaces (Workgroup, Private, etc) to LLVM address
diff --git a/mlir/include/mlir/Conversion/ArithToAMDGPU/ArithToAMDGPU.h b/mlir/include/mlir/Conversion/ArithToAMDGPU/ArithToAMDGPU.h
index fd144edf77452..e6e137759a76a 100644
--- a/mlir/include/mlir/Conversion/ArithToAMDGPU/ArithToAMDGPU.h
+++ b/mlir/include/mlir/Conversion/ArithToAMDGPU/ArithToAMDGPU.h
@@ -9,7 +9,7 @@
#ifndef MLIR_CONVERSION_ARITHTOAMDGPU_ARITHTOAMDGPU_H
#define MLIR_CONVERSION_ARITHTOAMDGPU_ARITHTOAMDGPU_H
-#include "mlir/Dialect/AMDGPU/Utils/Chipset.h"
+#include "mlir/Dialect/LLVMIR/ROCDLTargetInfo.h"
#include "mlir/IR/PatternMatch.h"
#include <memory>
#include <string>
@@ -31,7 +31,7 @@ namespace arith {
void populateArithToAMDGPUConversionPatterns(
RewritePatternSet &patterns, bool convertFP8Arithmetic,
bool saturateFP8Truncf, bool allowPackedF16Rtz, bool supportsScaledExtTrunc,
- amdgpu::Chipset chipset, PatternBenefit benefit = 1);
+ const ROCDL::TargetInfo &target, PatternBenefit benefit = 1);
} // namespace arith
} // namespace mlir
diff --git a/mlir/include/mlir/Conversion/GPUToROCDL/GPUToROCDLPass.h b/mlir/include/mlir/Conversion/GPUToROCDL/GPUToROCDLPass.h
index 220da0ad3c08f..494ee8cfa11ab 100644
--- a/mlir/include/mlir/Conversion/GPUToROCDL/GPUToROCDLPass.h
+++ b/mlir/include/mlir/Conversion/GPUToROCDL/GPUToROCDLPass.h
@@ -10,6 +10,7 @@
#include "mlir/Conversion/GPUToROCDL/Runtimes.h"
#include "mlir/Conversion/LLVMCommon/LoweringOptions.h"
+#include "mlir/Dialect/LLVMIR/ROCDLTargetInfo.h"
#include <memory>
namespace mlir {
@@ -21,10 +22,6 @@ class RewritePatternSet;
template <typename OpT>
class OperationPass;
-namespace amdgpu {
-struct Chipset;
-} // namespace amdgpu
-
namespace gpu {
class GPUModuleOp;
} // namespace gpu
@@ -38,7 +35,7 @@ class GPUModuleOp;
void populateGpuToROCDLConversionPatterns(const LLVMTypeConverter &converter,
RewritePatternSet &patterns,
gpu::amd::Runtime runtime,
- amdgpu::Chipset chipset);
+ const ROCDL::TargetInfo &target);
/// Configure target to convert from the GPU dialect to ROCDL.
void configureGpuToROCDLConversionLegality(ConversionTarget &target);
diff --git a/mlir/include/mlir/Conversion/MathToROCDL/MathToROCDL.h b/mlir/include/mlir/Conversion/MathToROCDL/MathToROCDL.h
index 60f1888569362..8ba104972abff 100644
--- a/mlir/include/mlir/Conversion/MathToROCDL/MathToROCDL.h
+++ b/mlir/include/mlir/Conversion/MathToROCDL/MathToROCDL.h
@@ -9,7 +9,7 @@
#define MLIR_CONVERSION_MATHTOROCDL_MATHTOROCDL_H_
#include "mlir/Conversion/LLVMCommon/TypeConverter.h"
-#include "mlir/Dialect/AMDGPU/Utils/Chipset.h"
+#include "mlir/Dialect/LLVMIR/ROCDLTargetInfo.h"
#include "mlir/IR/PatternMatch.h"
#include <memory>
@@ -20,11 +20,11 @@ class Pass;
#include "mlir/Conversion/Passes.h.inc"
/// Populate the given list with patterns that convert from Math to ROCDL calls.
-// `chipset` specifies the AMDGPU chipset to target. If `std::nullopt`,
-// none of the chipset dependent patterns are added.
+// `target` describes the AMDGPU target. If `std::nullopt`, none of the
+// target-dependent patterns are added.
void populateMathToROCDLConversionPatterns(
const LLVMTypeConverter &converter, RewritePatternSet &patterns,
- std::optional<amdgpu::Chipset> chipset);
+ std::optional<ROCDL::TargetInfo> target);
} // namespace mlir
#endif // MLIR_CONVERSION_MATHTOROCDL_MATHTOROCDL_H_
diff --git a/mlir/include/mlir/Conversion/Passes.td b/mlir/include/mlir/Conversion/Passes.td
index 6f3e288d454fe..2f80414d70b1f 100644
--- a/mlir/include/mlir/Conversion/Passes.td
+++ b/mlir/include/mlir/Conversion/Passes.td
@@ -130,9 +130,13 @@ def ConvertAMDGPUToROCDLPass : Pass<"convert-amdgpu-to-rocdl"> {
"LLVM::LLVMDialect",
"ROCDL::ROCDLDialect",
];
- let options = [Option<"chipset", "chipset", "std::string",
- /*default=*/"\"gfx000\"",
- "Chipset that these operations will run on">];
+ let options = [
+ Option<"arch", "arch", "std::string",
+ /*default=*/"\"invalid\"",
+ "Target architecture, as in Clang, with optional target-ID modifiers. New-style triples such as amdgpu9.42-amd-amdhsa are preferred, and can be extended to a full target ID like amdgpu9.42-amd-amdhsa--gfx942:xnack+:sramecc-. Bare chip names like gfx1250 or gfx942:xnack- are also supported. Defaults to an invalid target so that one must be given explicitly">,
+ Option<"chipset", "chipset", "std::string", /*default=*/"\"\"",
+ "Deprecated alias for 'arch'.">,
+ ];
}
//===----------------------------------------------------------------------===//
@@ -150,9 +154,11 @@ def ArithToAMDGPUConversionPass : Pass<"convert-arith-to-amdgpu"> {
let dependentDialects = ["amdgpu::AMDGPUDialect", "vector::VectorDialect"];
let options = [
- Option<"chipset", "chipset", "std::string",
- /*default=*/"\"gfx000\"",
- "Chipset that these operations will run on">,
+ Option<"arch", "arch", "std::string",
+ /*default=*/"\"invalid\"",
+ "Target architecture, as in Clang, with optional target-ID modifiers. New-style triples such as amdgpu9.42-amd-amdhsa are preferred, and can be extended to a full target ID like amdgpu9.42-amd-amdhsa--gfx942:xnack+:sramecc-. Bare chip names like gfx1250 or gfx942:xnack- are also supported. Defaults to an invalid target so that one must be given explicitly">,
+ Option<"chipset", "chipset", "std::string", /*default=*/"\"\"",
+ "Deprecated alias for 'arch'.">,
Option<"saturateFP8Truncf", "saturate-fp8-truncf", "bool",
/*default=*/"false",
"Use saturating truncation for 8-bit float types">,
@@ -684,9 +690,14 @@ def ConvertGpuOpsToROCDLOps : Pass<"convert-gpu-to-rocdl", "gpu::GPUModuleOp"> {
"memref::MemRefDialect",
];
let options = [
- Option<"chipset", "chipset", "std::string",
- /*default=*/"\"gfx000\"",
- "Chipset that these operations will run on">,
+ Option<"arch", "arch", "std::string",
+ /*default=*/"\"invalid\"",
+ "Target architecture, as in Clang, with optional target-ID modifiers. New-style triples such as amdgpu9.42-amd-amdhsa are preferred, and can be extended to a full target ID like amdgpu9.42-amd-amdhsa--gfx942:xnack+:sramecc-. Bare chip names like gfx1250 or gfx942:xnack- are also supported. Defaults to an invalid target so that one must be given explicitly">,
+ Option<"chipset", "chipset", "std::string", /*default=*/"\"\"",
+ "Deprecated alias for 'arch'.">,
+ Option<"waveSize", "wavesize", "unsigned", /*default=*/"0",
+ "Wavefront size (32 or 64) for targets that run at either, or 0 to "
+ "use the architecture's default">,
Option<"indexBitwidth", "index-bitwidth", "unsigned",
/*default=kDeriveIndexBitwidthFromDataLayout*/ "0",
"Bitwidth of the index type, 0 to use size of machine word">,
@@ -851,9 +862,9 @@ def ConvertMathToROCDL : Pass<"convert-math-to-rocdl", "ModuleOp"> {
let description = [{
This pass converts supported Math ops to ROCDL library calls.
- The chipset option specifies the target AMDGPU architecture. If the chipset
- is empty, none of the chipset-dependent patterns are added, and the pass
- will not attempt to parse the chipset.
+ The triple option specifies the target AMDGPU architecture. If it is empty,
+ none of the target-dependent patterns are added and the pass does not
+ resolve a target; a chip or feature list without a triple is rejected.
}];
let dependentDialects = [
"arith::ArithDialect",
@@ -861,9 +872,14 @@ def ConvertMathToROCDL : Pass<"convert-math-to-rocdl", "ModuleOp"> {
"ROCDL::ROCDLDialect",
"vector::VectorDialect",
];
- let options = [Option<"chipset", "chipset", "std::string",
- /*default=*/"\"\"",
- "Chipset that these operations will run on">];
+ let options = [
+ Option<"arch", "arch", "std::string", /*default=*/"\"\"",
+ "Target architecture, as in Clang, with optional target-ID modifiers "
+ "(e.g. amdgpu9.42-amd-amdhsa, gfx942, gfx942:xnack-). "
+ "If empty, no target-dependent patterns are added">,
+ Option<"chipset", "chipset", "std::string", /*default=*/"\"\"",
+ "Deprecated alias for 'arch'.">,
+ ];
}
//===----------------------------------------------------------------------===//
diff --git a/mlir/include/mlir/Dialect/AMDGPU/Transforms/Passes.h b/mlir/include/mlir/Dialect/AMDGPU/Transforms/Passes.h
index 48e7658568f86..c8ec368231f83 100644
--- a/mlir/include/mlir/Dialect/AMDGPU/Transforms/Passes.h
+++ b/mlir/include/mlir/Dialect/AMDGPU/Transforms/Passes.h
@@ -13,7 +13,7 @@
#ifndef MLIR_DIALECT_AMDGPU_TRANSFORMS_PASSES_H_
#define MLIR_DIALECT_AMDGPU_TRANSFORMS_PASSES_H_
-#include "mlir/Dialect/AMDGPU/Utils/Chipset.h"
+#include "mlir/Dialect/LLVMIR/ROCDLTargetInfo.h"
#include "mlir/IR/PatternMatch.h"
#include "mlir/Pass/Pass.h"
@@ -29,7 +29,7 @@ namespace amdgpu {
void populateAmdgpuEmulateAtomicsPatterns(ConversionTarget &target,
RewritePatternSet &patterns,
- Chipset chipset,
+ const ROCDL::TargetInfo &targetInfo,
PatternBenefit benefit = 1);
void populateAmdgpuResolveStridedMetadataPatterns(RewritePatternSet &patterns,
diff --git a/mlir/include/mlir/Dialect/AMDGPU/Transforms/Passes.td b/mlir/include/mlir/Dialect/AMDGPU/Transforms/Passes.td
index 7dd7ac750a9eb..d9e68a7be4dc7 100644
--- a/mlir/include/mlir/Dialect/AMDGPU/Transforms/Passes.td
+++ b/mlir/include/mlir/Dialect/AMDGPU/Transforms/Passes.td
@@ -16,19 +16,23 @@
include "mlir/Pass/PassBase.td"
def AmdgpuEmulateAtomicsPass : Pass<"amdgpu-emulate-atomics"> {
- let summary = "Emulate atomic operations on chipsets that do not support them";
+ let summary = "Emulate atomic operations the target does not support";
let description = [{
- This pass rewrites any AMDGPU-specific atomic operation that is not supported
- on the given `chipset` into a compare-and-swap loop.
+ This pass rewrites any AMDGPU-specific atomic operation that the target does
+ not support into a compare-and-swap loop.
}];
let dependentDialects = [
"cf::ControlFlowDialect",
"arith::ArithDialect",
"vector::VectorDialect"
];
- let options = [Option<"chipset", "chipset", "std::string",
- /*default=*/"\"gfx000\"",
- "Chipset that these operations will run on">];
+ let options = [
+ Option<"arch", "arch", "std::string",
+ /*default=*/"\"invalid\"",
+ "Target architecture, as in Clang, with optional target-ID modifiers. New-style triples such as amdgpu9.42-amd-amdhsa are preferred, and can be extended to a full target ID like amdgpu9.42-amd-amdhsa--gfx942:xnack+:sramecc-. Bare chip names like gfx1250 or gfx942:xnack- are also supported. Defaults to an invalid target so that one must be given explicitly">,
+ Option<"chipset", "chipset", "std::string", /*default=*/"\"\"",
+ "Deprecated alias for 'arch'.">,
+ ];
}
def AmdgpuResolveStridedMetadataPass : Pass<"amdgpu-resolve-strided-metadata"> {
diff --git a/mlir/include/mlir/Dialect/AMDGPU/Utils/Chipset.h b/mlir/include/mlir/Dialect/AMDGPU/Utils/Chipset.h
index 256065537aeb9..1a964c27a6333 100644
--- a/mlir/include/mlir/Dialect/AMDGPU/Utils/Chipset.h
+++ b/mlir/include/mlir/Dialect/AMDGPU/Utils/Chipset.h
@@ -9,6 +9,7 @@
#define MLIR_DIALECT_AMDGPU_UTILS_CHIPSET_H_
#include "mlir/Support/LLVM.h"
+#include "llvm/Support/Compiler.h"
#include <tuple>
namespace mlir::amdgpu {
@@ -19,6 +20,10 @@ namespace mlir::amdgpu {
/// gfx942 --> major = 9, minor = 0x4, stepping = 0x2
/// gfx90a --> major = 9, minor = 0x0, stepping = 0xa
/// gfx1103 --> major = 11, minor = 0x0, stepping = 0x3
+///
+/// \deprecated Use `mlir::ROCDL::TargetInfo` instead, and rely on target
+/// features rather than about version numbers. Will be removed after one
+/// release.
struct Chipset {
unsigned majorVersion = 0; // The major version (decimal).
unsigned minorVersion = 0; // The minor version (hexadecimal).
@@ -30,6 +35,9 @@ struct Chipset {
/// Parses the chipset version string and returns the chipset on success, and
/// failure otherwise.
+ ///
+ /// \deprecated Use `ROCDL::TargetInfo::get`.
+ LLVM_DEPRECATED("use ROCDL::TargetInfo::get instead", "")
static FailureOr<Chipset> parse(StringRef name);
std::tuple<unsigned, unsigned, unsigned> asTuple() const {
@@ -49,6 +57,10 @@ struct Chipset {
#undef DEFINE_COMP_OPERATOR
};
+/// \deprecated Test `llvm::AMDGPU::FEAT_OCP_FP8_CONVERSION_INSTS` on a
+/// `ROCDL::TargetInfo` instead. This misses gfx11.7, which does have the OCP
+/// fp8 conversions.
+LLVM_DEPRECATED("test FEAT_OCP_FP8_CONVERSION_INSTS on a ROCDL::TargetInfo", "")
inline bool hasOcpFp8(const Chipset &chipset) {
return (chipset.majorVersion == 9 && chipset.minorVersion >= 5) ||
chipset.majorVersion >= 12;
diff --git a/mlir/include/mlir/Dialect/GPU/Pipelines/Passes.h b/mlir/include/mlir/Dialect/GPU/Pipelines/Passes.h
index 6327415d62769..c34cc0ee9fdd6 100644
--- a/mlir/include/mlir/Dialect/GPU/Pipelines/Passes.h
+++ b/mlir/include/mlir/Dialect/GPU/Pipelines/Passes.h
@@ -72,17 +72,16 @@ struct GPUToROCDLPipelineOptions
llvm::cl::desc("Bitwidth of the index type for the host (warning this "
"should be 64 until the GPU layering is fixed)"),
llvm::cl::init(64)};
- PassOptions::Option<std::string> triple{
- *this, "triple",
- llvm::cl::desc("AMDGPU target triple (e.g. amdgcn-amd-amdhsa)."),
- llvm::cl::init("amdgcn-amd-amdhsa")};
- PassOptions::Option<std::string> chip{
- *this, "chip",
+ PassOptions::Option<std::string> arch{
+ *this, "arch",
llvm::cl::desc(
- "AMDGPU target chip (e.g. gfx90a, gfx942, gfx1100). Required: "
- "AMDGCN binaries are not forward-compatible across chip families.")};
- PassOptions::Option<std::string> features{
- *this, "features", llvm::cl::desc("AMDGPU target features."),
+ "AMDGPU target architecture, as in Clang, with optional target-ID "
+ "modifiers (e.g. gfx942, gfx90a:xnack+, "
+ "amdgpu9.0a-amd-amdhsa--gfx90a:xnack-). Required: AMDGCN binaries "
+ "are "
+ "not forward-compatible across chip families.")};
+ PassOptions::Option<std::string> chip{
+ *this, "chip", llvm::cl::desc("Deprecated alias for 'arch'."),
llvm::cl::init("")};
PassOptions::Option<std::string> binaryFormat{
*this, "binary-format",
@@ -93,11 +92,11 @@ struct GPUToROCDLPipelineOptions
*this, "abi",
llvm::cl::desc("AMDHSA ABI version (e.g. \"500\", \"600\")."),
llvm::cl::init("600")};
- PassOptions::Option<bool> wave64{
- *this, "wave64",
- llvm::cl::desc("Use Wave64 mode (default true; wave32 if false, "
- "appropriate for RDNA / gfx10+ where supported)."),
- llvm::cl::init(true)};
+ PassOptions::Option<unsigned> waveSize{
+ *this, "wavesize",
+ llvm::cl::desc("Wavefront size (32 or 64) for targets that run at "
+ "either, or 0 to use the architecture's default."),
+ llvm::cl::init(0)};
PassOptions::Option<int> optLevel{
*this, "opt-level",
llvm::cl::desc("Optimization level for ROCDL/AMDGPU compilation."),
diff --git a/mlir/include/mlir/Dialect/GPU/TransformOps/GPUTransformOps.td b/mlir/include/mlir/Dialect/GPU/TransformOps/GPUTransformOps.td
index 030506bba1076..d56cc1f13afd2 100644
--- a/mlir/include/mlir/Dialect/GPU/TransformOps/GPUTransformOps.td
+++ b/mlir/include/mlir/Dialect/GPU/TransformOps/GPUTransformOps.td
@@ -61,11 +61,18 @@ def ApplyGPUToROCDLConversionPatternsOp : Op<Transform_Dialect,
let description = [{
Collects patterns that convert GPU dialect ops to ROCDL dialect ops. These
patterns require an "LLVMTypeConverter".
+
+ `arch` names the target the way Clang does: a GPU name with optional
+ target-ID modifiers ("gfx942", "gfx942:xnack+", "gfx9-4-generic"), a triple
+ ("amdgpu9.42-amd-amdhsa"), or a full target ID
+ ("amdgpu9.0a-amd-amdhsa--gfx90a:sramecc+:xnack-").
+
+ `wavesize` pins the wavefront size on targets that run at either; omitting
+ it uses the architecture's own default.
}];
- let arguments = (ins StrAttr:$chipset);
- let assemblyFormat = [{
- `chipset` `=` $chipset attr-dict
- }];
+ let arguments = (ins StrAttr:$arch,
+ OptionalAttr<I32Attr>:$wavesize);
+ let assemblyFormat = "prop-dict attr-dict";
}
//===----------------------------------------------------------------------===//
@@ -328,11 +335,15 @@ def ApplyGPUPromoteShuffleToAMDGPUPatternsOp : Op<Transform_Dialect,
let description = [{
Collects patterns that are tryin to promote `gpu.shuffle`s to specialized
AMDGPU intrinsics.
+
+ `arch` names the target the way Clang does: a GPU name with optional
+ target-ID modifiers ("gfx950", "gfx950:xnack+"), a triple
+ ("amdgpu9.50-amd-amdhsa"), or a full target ID. Omitting it collects only
+ the patterns that hold on every target.
}];
- let arguments = (ins OptionalAttr<StrAttr>:$chipset);
- let assemblyFormat = [{
- (`chipset` `=` $chipset^)? attr-dict
- }];
+ let arguments = (ins OptionalAttr<StrAttr>:$arch);
+ let assemblyFormat = "prop-dict attr-dict";
+ let hasVerifier = 1;
}
diff --git a/mlir/include/mlir/Dialect/GPU/Transforms/Passes.h b/mlir/include/mlir/Dialect/GPU/Transforms/Passes.h
index d5c253d6c9c08..6657ff819dea5 100644
--- a/mlir/include/mlir/Dialect/GPU/Transforms/Passes.h
+++ b/mlir/include/mlir/Dialect/GPU/Transforms/Passes.h
@@ -13,9 +13,9 @@
#ifndef MLIR_DIALECT_GPU_TRANSFORMS_PASSES_H_
#define MLIR_DIALECT_GPU_TRANSFORMS_PASSES_H_
-#include "mlir/Dialect/AMDGPU/Utils/Chipset.h"
#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/GPU/Utils/GPUUtils.h"
+#include "mlir/Dialect/LLVMIR/ROCDLTargetInfo.h"
#include "mlir/IR/PatternMatch.h"
#include "mlir/Pass/Pass.h"
#include <optional>
@@ -76,16 +76,15 @@ void populateGpuLowerClusteredSubgroupReduceToShufflePatterns(
/// Collect a set of patterns to lower `gpu.subgroup_reduce` into `amdgpu.dpp`
/// ops over scalar types. Assumes that the subgroup has
/// `subgroupSize` lanes. Applicable only to AMD GPUs.
-void populateGpuLowerSubgroupReduceToDPPPatterns(RewritePatternSet &patterns,
- unsigned subgroupSize,
- amdgpu::Chipset chipset,
- PatternBenefit benefit = 1);
+void populateGpuLowerSubgroupReduceToDPPPatterns(
+ RewritePatternSet &patterns, unsigned subgroupSize,
+ const ROCDL::TargetInfo &target, PatternBenefit benefit = 1);
/// Disjoint counterpart of `populateGpuLowerSubgroupReduceToDPPPatterns`
/// that only matches `gpu.subgroup_reduce` ops with a `cluster_size`.
void populateGpuLowerClusteredSubgroupReduceToDPPPatterns(
- RewritePatternSet &patterns, unsigned subgroupSize, amdgpu::Chipset chipset,
- PatternBenefit benefit = 1);
+ RewritePatternSet &patterns, unsigned subgroupSize,
+ const ROCDL::TargetInfo &target, PatternBenefit benefit = 1);
/// Collect all patterns to rewrite ops within the GPU dialect.
inline void populateGpuRewritePatterns(RewritePatternSet &patterns) {
@@ -115,7 +114,7 @@ void populateGpuEliminateBarriersPatterns(RewritePatternSet &patterns);
/// Tries to promote `gpu.shuffle`s to specialized AMDGPU intrinsics.
void populateGpuPromoteShuffleToAMDGPUPatterns(
- RewritePatternSet &patterns, std::optional<amdgpu::Chipset> maybeChipset);
+ RewritePatternSet &patterns, std::optional<ROCDL::TargetInfo> target);
/// Generate the code for registering passes.
#define GEN_PASS_REGISTRATION
diff --git a/mlir/include/mlir/Dialect/GPU/Transforms/Passes.td b/mlir/include/mlir/Dialect/GPU/Transforms/Passes.td
index 1aaab1ee91523..090bfdba3ce8f 100644
--- a/mlir/include/mlir/Dialect/GPU/Transforms/Passes.td
+++ b/mlir/include/mlir/Dialect/GPU/Transforms/Passes.td
@@ -191,12 +191,20 @@ def GpuROCDLAttachTarget: Pass<"rocdl-attach-target", ""> {
Option<"moduleMatcher", "module", "std::string",
/*default=*/ [{""}],
"Regex used to identify the modules to attach the target to.">,
+ Option<"arch", "arch", "std::string",
+ /*default=*/"\"\"",
+ "Target architecture, in the spelling the conversion passes take "
+ "(e.g. gfx942, gfx90a:xnack+, amdgpu9.4-amd-amdhsa, "
+ "amdgcn-amd-amdhsa--gfx90a:xnack-). "
+ "When non-empty this supersedes 'triple' and 'chip', and its "
+ "target-ID modifiers become the module's 'rocdl.xnack' and "
+ "'rocdl.sramecc' attributes rather than target features.">,
Option<"triple", "triple", "std::string",
/*default=*/ "\"amdgcn-amd-amdhsa\"",
- "Target triple.">,
+ "Target triple. Ignored when 'arch' is given.">,
Option<"chip", "chip", "std::string",
/*default=*/"\"gfx900\"",
- "Target chip.">,
+ "Target chip. Ignored when 'arch' is given.">,
Option<"features", "features", "std::string",
/*default=*/"\"\"",
"Target features.">,
@@ -206,9 +214,10 @@ def GpuROCDLAttachTarget: Pass<"rocdl-attach-target", ""> {
Option<"optLevel", "O", "unsigned",
/*default=*/"2",
"Optimization level.">,
- Option<"wave64Flag", "wave64", "bool",
- /*default=*/"true",
- "Use Wave64 mode.">,
+ Option<"waveSize", "wavesize", "unsigned",
+ /*default=*/"0",
+ "Wavefront size (32 or 64). 0 takes the size from the target when "
+ "'arch' is given, and Wave64 otherwise.">,
Option<"fastFlag", "fast", "bool",
/*default=*/"false",
"Enable fast relaxed math opt.">,
diff --git a/mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp b/mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp
index 6bbe65731964f..4cb739816e97d 100644
--- a/mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp
+++ b/mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp
@@ -14,7 +14,6 @@
#include "mlir/Conversion/LLVMCommon/TypeConverter.h"
#include "mlir/Dialect/AMDGPU/IR/AMDGPUDialect.h"
#include "mlir/Dialect/AMDGPU/IR/AMDGPUEnums.h"
-#include "mlir/Dialect/AMDGPU/Utils/Chipset.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
#include "mlir/Dialect/LLVMIR/LLVMTypes.h"
#include "mlir/Dialect/LLVMIR/ROCDLDialect.h"
@@ -32,6 +31,7 @@
#include "llvm/Support/AMDGPUAddrSpace.h"
#include "llvm/Support/Casting.h"
#include "llvm/Support/ErrorHandling.h"
+#include "llvm/Support/MathExtras.h"
#include <cstdint>
#include <optional>
@@ -43,72 +43,6 @@ namespace mlir {
using namespace mlir;
using namespace mlir::amdgpu;
-// Define commonly used chipsets versions for convenience.
-constexpr Chipset kGfx908 = Chipset(9, 0, 8);
-constexpr Chipset kGfx90a = Chipset(9, 0, 0xa);
-constexpr Chipset kGfx942 = Chipset(9, 4, 2);
-constexpr Chipset kGfx950 = Chipset(9, 5, 0);
-constexpr Chipset kGfx1200 = Chipset(12, 0, 0);
-constexpr Chipset kGfx1250 = Chipset(12, 5, 0);
-
-// Predicates mirroring the LLVM AMDGPU `HasDot{N}Insts` features that gate
-// the `v_dot*` instructions consumed by the `amdgpu.dot` lowering.
-static bool hasDot1Insts(const Chipset &chipset) {
- if (chipset.majorVersion == 9)
- return chipset >= Chipset(9, 0, 6);
- if (chipset.majorVersion == 10) {
- if (chipset.minorVersion == 1)
- return chipset.steppingVersion == 1u || chipset.steppingVersion == 2u;
- return chipset.minorVersion >= 3u;
- }
- return false;
-}
-
-static bool hasDot2Insts(const Chipset &chipset) {
- return hasDot1Insts(chipset);
-}
-
-static bool hasDot7Insts(const Chipset &chipset) {
- return chipset.majorVersion >= 11 || hasDot1Insts(chipset);
-}
-
-static bool hasDot8Insts(const Chipset &chipset) {
- return chipset.majorVersion >= 11;
-}
-
-static bool hasDot9Insts(const Chipset &chipset) {
- if (chipset.majorVersion == 11)
- return true;
- return chipset.majorVersion == 12 && chipset.minorVersion == 0;
-}
-
-static bool hasDot10Insts(const Chipset &chipset) {
- if (chipset.majorVersion == 11)
- return true;
- if (chipset.majorVersion == 12)
- return chipset.minorVersion == 0;
- return hasDot1Insts(chipset);
-}
-
-static bool hasDot11Insts(const Chipset &chipset) {
- if (chipset.majorVersion == 11)
- return chipset.minorVersion == 7u;
- return chipset.majorVersion == 12 && chipset.minorVersion == 0;
-}
-
-static bool hasDot12Insts(const Chipset &chipset) {
- if (chipset == Chipset(9, 5, 0))
- return true;
- if (chipset.majorVersion == 11)
- return true;
- return chipset.majorVersion == 12 && chipset.minorVersion == 0;
-}
-
-static bool has45BitNumRecordsBufferResource(const Chipset &chipset) {
- return chipset.majorVersion > 12 ||
- (chipset.majorVersion == 12 && chipset.minorVersion >= 5);
-}
-
/// Zero-extend or truncate the unsigned number `val` to `width` bits.
static Value convertUnsignedToInt(ConversionPatternRewriter &rewriter,
Location loc, Value val, unsigned width) {
@@ -172,10 +106,9 @@ static Value getNumRecords(ConversionPatternRewriter &rewriter, Location loc,
MemRefType memrefType,
MemRefDescriptor &memrefDescriptor,
ArrayRef<int64_t> strides, int64_t elementByteWidth,
- amdgpu::Chipset chipset, bool boundsCheck) {
- if (has45BitNumRecordsBufferResource(chipset) && !boundsCheck) {
- constexpr int64_t first45bits = (1ll << 45) - 1;
- return createI64Constant(rewriter, loc, first45bits);
+ unsigned numRecordsWidth, bool boundsCheck) {
+ if (numRecordsWidth > 32 && !boundsCheck) {
+ return createI64Constant(rewriter, loc, llvm::maxUIntN(numRecordsWidth));
}
if (memrefType.hasStaticShape() &&
!llvm::any_of(strides, ShapedType::isDynamic)) {
@@ -202,7 +135,8 @@ static Value getNumRecords(ConversionPatternRewriter &rewriter, Location loc,
static Value makeBufferRsrc(ConversionPatternRewriter &rewriter, Location loc,
Value basePointer, Value numRecords,
- bool boundsCheck, amdgpu::Chipset chipset,
+ bool boundsCheck, const ROCDL::TargetInfo &target,
+ unsigned numRecordsWidth,
Value cacheSwizzleStride = nullptr,
unsigned addressSpace = 8) {
// The stride value is generally 0. However, on MI-300 and onward, you can
@@ -210,7 +144,7 @@ static Value makeBufferRsrc(ConversionPatternRewriter &rewriter, Location loc,
// and setting that stride to a cache stride.
Type i16 = rewriter.getI16Type();
Value stride;
- if (chipset.majorVersion == 9 && chipset >= kGfx942 && cacheSwizzleStride) {
+ if (target.has(llvm::AMDGPU::FEAT_GFX940_INSTS) && cacheSwizzleStride) {
Value cacheStrideZext =
LLVM::ZExtOp::create(rewriter, loc, i16, cacheSwizzleStride);
Value swizzleBit = LLVM::ConstantOp::create(
@@ -223,7 +157,7 @@ static Value makeBufferRsrc(ConversionPatternRewriter &rewriter, Location loc,
}
uint32_t flags = 0;
- if (chipset >= kGfx1250) {
+ if (target.has(llvm::AMDGPU::FEAT_GFX1250_INSTS)) {
// Flag word:
// bit 0: swizzle
// bit 1: 0 means (total_offset + payload > numRecords)
@@ -249,16 +183,14 @@ static Value makeBufferRsrc(ConversionPatternRewriter &rewriter, Location loc,
// none, 3 = either swizzles or testing against offset field) RDNA only
// bits 30-31: Type (must be 0)
flags |= (7 << 12) | (4 << 15);
- if (chipset.majorVersion >= 10) {
+ if (target.has(llvm::AMDGPU::FEAT_GFX10_INSTS)) {
flags |= (1 << 24);
uint32_t oob = boundsCheck ? 3 : 2;
flags |= (oob << 28);
}
}
Value flagsConst = createI32Constant(rewriter, loc, flags);
- numRecords =
- convertUnsignedToInt(rewriter, loc, numRecords,
- has45BitNumRecordsBufferResource(chipset) ? 45 : 32);
+ numRecords = convertUnsignedToInt(rewriter, loc, numRecords, numRecordsWidth);
Type rsrcType =
LLVM::LLVMPointerType::get(rewriter.getContext(), addressSpace);
Value resource = rewriter.createOrFold<ROCDL::MakeBufferRsrcOp>(
@@ -269,11 +201,11 @@ static Value makeBufferRsrc(ConversionPatternRewriter &rewriter, Location loc,
namespace {
struct FatRawBufferCastLowering
: public ConvertOpToLLVMPattern<FatRawBufferCastOp> {
- FatRawBufferCastLowering(const LLVMTypeConverter &converter, Chipset chipset)
- : ConvertOpToLLVMPattern<FatRawBufferCastOp>(converter),
- chipset(chipset) {}
+ FatRawBufferCastLowering(const LLVMTypeConverter &converter,
+ const ROCDL::TargetInfo &target)
+ : ConvertOpToLLVMPattern<FatRawBufferCastOp>(converter), target(target) {}
- Chipset chipset;
+ ROCDL::TargetInfo target;
LogicalResult
matchAndRewrite(FatRawBufferCastOp op, FatRawBufferCastOpAdaptor adaptor,
@@ -293,11 +225,17 @@ struct FatRawBufferCastLowering
if (failed(memrefType.getStridesAndOffset(strideVals, unusedOffset)))
return op.emitOpError("Can't lower non-stride-offset memrefs");
+ std::optional<unsigned> numRecordsWidth =
+ target.getBufferResourceNumRecordsWidth();
+ if (!numRecordsWidth)
+ return op.emitOpError(
+ "buffer resource num_records width is unknown for this target");
+
Value numRecords = adaptor.getValidBytes();
if (!numRecords)
- numRecords =
- getNumRecords(rewriter, loc, memrefType, descriptor, strideVals,
- elementByteWidth, chipset, adaptor.getBoundsCheck());
+ numRecords = getNumRecords(rewriter, loc, memrefType, descriptor,
+ strideVals, elementByteWidth, *numRecordsWidth,
+ adaptor.getBoundsCheck());
Value basePointer =
adaptor.getResetOffset()
@@ -324,7 +262,8 @@ struct FatRawBufferCastLowering
Value fatPtr = makeBufferRsrc(
rewriter, loc, basePointer, numRecords, adaptor.getBoundsCheck(),
- chipset, adaptor.getCacheSwizzleStride(), /*addressSpace=*/7);
+ target, *numRecordsWidth, adaptor.getCacheSwizzleStride(),
+ /*addressSpace=*/7);
Value result = MemRefDescriptor::poison(
rewriter, loc,
@@ -349,10 +288,11 @@ struct FatRawBufferCastLowering
/// Define lowering patterns for raw buffer ops
template <typename GpuOp, typename Intrinsic>
struct RawBufferOpLowering : public ConvertOpToLLVMPattern<GpuOp> {
- RawBufferOpLowering(const LLVMTypeConverter &converter, Chipset chipset)
- : ConvertOpToLLVMPattern<GpuOp>(converter), chipset(chipset) {}
+ RawBufferOpLowering(const LLVMTypeConverter &converter,
+ const ROCDL::TargetInfo &target)
+ : ConvertOpToLLVMPattern<GpuOp>(converter), target(target) {}
- Chipset chipset;
+ ROCDL::TargetInfo target;
static constexpr uint32_t maxVectorOpWidth = 128;
LogicalResult
@@ -363,7 +303,7 @@ struct RawBufferOpLowering : public ConvertOpToLLVMPattern<GpuOp> {
Value unconvertedMemref = gpuOp.getMemref();
MemRefType memrefType = cast<MemRefType>(unconvertedMemref.getType());
- if (chipset.majorVersion < 9)
+ if (!target.has(llvm::AMDGPU::FEAT_GFX9_INSTS))
return gpuOp.emitOpError("raw buffer ops require GCN or higher");
Value storeData = adaptor.getODSOperands(0)[0];
@@ -468,11 +408,18 @@ struct RawBufferOpLowering : public ConvertOpToLLVMPattern<GpuOp> {
Value ptr = memrefDescriptor.bufferPtr(
rewriter, loc, *this->getTypeConverter(), memrefType);
- Value numRecords =
- getNumRecords(rewriter, loc, memrefType, memrefDescriptor, strides,
- elementByteWidth, chipset, adaptor.getBoundsCheck());
- Value resource = makeBufferRsrc(rewriter, loc, ptr, numRecords,
- adaptor.getBoundsCheck(), chipset);
+ std::optional<unsigned> numRecordsWidth =
+ target.getBufferResourceNumRecordsWidth();
+ if (!numRecordsWidth)
+ return gpuOp.emitOpError(
+ "buffer resource num_records width is unknown for this target");
+
+ Value numRecords = getNumRecords(
+ rewriter, loc, memrefType, memrefDescriptor, strides, elementByteWidth,
+ *numRecordsWidth, adaptor.getBoundsCheck());
+ Value resource =
+ makeBufferRsrc(rewriter, loc, ptr, numRecords, adaptor.getBoundsCheck(),
+ target, *numRecordsWidth);
args.push_back(resource);
// Indexing (voffset)
@@ -526,15 +473,18 @@ struct RawBufferOpLowering : public ConvertOpToLLVMPattern<GpuOp> {
/// Lgkmcnt = Waitcnt[11:8] (pre-gfx10)
/// Lgkmcnt = Waitcnt[13:8] (gfx10)
/// Lgkmcnt = Waitcnt[9:4] (gfx11)
-static FailureOr<unsigned> encodeWaitcnt(Chipset chipset, unsigned vmcnt,
- unsigned expcnt, unsigned lgkmcnt) {
- if (chipset.majorVersion < 9) {
+static FailureOr<unsigned> encodeWaitcnt(const ROCDL::TargetInfo &target,
+ unsigned vmcnt, unsigned expcnt,
+ unsigned lgkmcnt) {
+ if (target.isUnknown())
+ return failure();
+ if (!target.has(llvm::AMDGPU::FEAT_GFX9_INSTS)) {
vmcnt = std::min(15u, vmcnt);
expcnt = std::min(7u, expcnt);
lgkmcnt = std::min(15u, lgkmcnt);
return vmcnt | (expcnt << 4) | (lgkmcnt << 8);
}
- if (chipset.majorVersion == 9) {
+ if (target.isGeneration(9)) {
vmcnt = std::min(63u, vmcnt);
expcnt = std::min(7u, expcnt);
lgkmcnt = std::min(15u, lgkmcnt);
@@ -543,7 +493,7 @@ static FailureOr<unsigned> encodeWaitcnt(Chipset chipset, unsigned vmcnt,
unsigned otherCnts = (expcnt << 4) | (lgkmcnt << 8);
return lowBits | highBits | otherCnts;
}
- if (chipset.majorVersion == 10) {
+ if (target.isGeneration(10)) {
vmcnt = std::min(63u, vmcnt);
expcnt = std::min(7u, expcnt);
lgkmcnt = std::min(63u, lgkmcnt);
@@ -552,7 +502,7 @@ static FailureOr<unsigned> encodeWaitcnt(Chipset chipset, unsigned vmcnt,
unsigned otherCnts = (expcnt << 4) | (lgkmcnt << 8);
return lowBits | highBits | otherCnts;
}
- if (chipset.majorVersion == 11) {
+ if (target.isGeneration(11)) {
vmcnt = std::min(63u, vmcnt);
expcnt = std::min(7u, expcnt);
lgkmcnt = std::min(63u, lgkmcnt);
@@ -564,16 +514,16 @@ static FailureOr<unsigned> encodeWaitcnt(Chipset chipset, unsigned vmcnt,
struct MemoryCounterWaitOpLowering
: public ConvertOpToLLVMPattern<MemoryCounterWaitOp> {
MemoryCounterWaitOpLowering(const LLVMTypeConverter &converter,
- Chipset chipset)
- : ConvertOpToLLVMPattern<MemoryCounterWaitOp>(converter),
- chipset(chipset) {}
+ const ROCDL::TargetInfo &target)
+ : ConvertOpToLLVMPattern<MemoryCounterWaitOp>(converter), target(target) {
+ }
- Chipset chipset;
+ ROCDL::TargetInfo target;
LogicalResult
matchAndRewrite(MemoryCounterWaitOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
- if (chipset.majorVersion >= 12) {
+ if (target.has(llvm::AMDGPU::FEAT_GFX12_INSTS)) {
Location loc = op.getLoc();
if (std::optional<int> ds = adaptor.getDs())
ROCDL::WaitDscntOp::create(rewriter, loc, *ds);
@@ -618,7 +568,7 @@ struct MemoryCounterWaitOpLowering
vmcnt = getVal(store);
}
- FailureOr<unsigned> waitcnt = encodeWaitcnt(chipset, vmcnt, exp, ds);
+ FailureOr<unsigned> waitcnt = encodeWaitcnt(target, vmcnt, exp, ds);
if (failed(waitcnt))
return op.emitOpError("unsupported chipset");
@@ -628,18 +578,24 @@ struct MemoryCounterWaitOpLowering
};
struct LDSBarrierOpLowering : public ConvertOpToLLVMPattern<LDSBarrierOp> {
- LDSBarrierOpLowering(const LLVMTypeConverter &converter, Chipset chipset)
- : ConvertOpToLLVMPattern<LDSBarrierOp>(converter), chipset(chipset) {}
+ LDSBarrierOpLowering(const LLVMTypeConverter &converter,
+ const ROCDL::TargetInfo &target)
+ : ConvertOpToLLVMPattern<LDSBarrierOp>(converter), target(target) {}
- Chipset chipset;
+ ROCDL::TargetInfo target;
LogicalResult
matchAndRewrite(LDSBarrierOp op, LDSBarrierOp::Adaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
Location loc = op.getLoc();
- // This ensures that waits on global memory aren't introduced on
- // chips that don't have the BackOffBarrier feature enabled in LLVM.
- bool requiresInlineAsm = chipset < kGfx90a;
+ bool hasSplitBarriers = target.has(llvm::AMDGPU::FEAT_GFX12_INSTS);
+ // Inline assembly is only needed for cases where the hardware doesn't use
+ // split barriers and doesn't have FeatureBackOffBarrier (this is mainly
+ // early gfx9). In that case, we use inline assembly to bypass the
+ // conservative insertion of global memory waits at barriers, since we care
+ // more about performance than having debug watches work correctly.
+ bool requiresInlineAsm =
+ !hasSplitBarriers && !target.has(llvm::AMDGPU::FEAT_BACK_OFF_BARRIER);
Attribute mmra =
rewriter.getAttr<LLVM::MMRATagAttr>("amdgpu-synchronize-as", "local");
@@ -669,7 +625,7 @@ struct LDSBarrierOpLowering : public ConvertOpToLLVMPattern<LDSBarrierOp> {
/*convergent=*/false,
/*asm_dialect=*/asmDialectAttr,
/*operand_attrs=*/ArrayAttr());
- } else if (chipset.majorVersion < 12) {
+ } else if (!hasSplitBarriers) {
ROCDL::SBarrierOp::create(rewriter, loc);
} else {
ROCDL::BarrierSignalOp::create(rewriter, loc, -1);
@@ -685,10 +641,11 @@ struct LDSBarrierOpLowering : public ConvertOpToLLVMPattern<LDSBarrierOp> {
};
struct SchedBarrierOpLowering : public ConvertOpToLLVMPattern<SchedBarrierOp> {
- SchedBarrierOpLowering(const LLVMTypeConverter &converter, Chipset chipset)
- : ConvertOpToLLVMPattern<SchedBarrierOp>(converter), chipset(chipset) {}
+ SchedBarrierOpLowering(const LLVMTypeConverter &converter,
+ const ROCDL::TargetInfo &target)
+ : ConvertOpToLLVMPattern<SchedBarrierOp>(converter), target(target) {}
- Chipset chipset;
+ ROCDL::TargetInfo target;
LogicalResult
matchAndRewrite(SchedBarrierOp op, SchedBarrierOp::Adaptor adaptor,
@@ -904,31 +861,34 @@ static void wmmaPushOutputOperand(ConversionPatternRewriter &rewriter,
}
/// Return true if `type` is the E5M2 variant of an 8-bit float that is
-/// supported by the `_bf8` instructions on the given `chipset`.
-static bool typeIsExpectedBf8ForChipset(Chipset chipset, Type type) {
- return (chipset == kGfx942 && isa<Float8E5M2FNUZType>(type)) ||
- (hasOcpFp8(chipset) && isa<Float8E5M2Type>(type));
+/// supported by the `_bf8` instructions on `target`.
+static bool typeIsExpectedBf8ForTarget(const ROCDL::TargetInfo &target,
+ Type type) {
+ return (target.hasFnuzFp8() && isa<Float8E5M2FNUZType>(type)) ||
+ (target.hasOcpFp8() && isa<Float8E5M2Type>(type));
}
/// Return true if `type` is the E4M3FN variant of an 8-bit float that is
-/// supported by the `_fp8` instructions on the given `chipset`.
-static bool typeIsExpectedFp8ForChipset(Chipset chipset, Type type) {
- return (chipset == kGfx942 && isa<Float8E4M3FNUZType>(type)) ||
- (hasOcpFp8(chipset) && isa<Float8E4M3FNType>(type));
+/// supported by the `_fp8` instructions on `target`.
+static bool typeIsExpectedFp8ForTarget(const ROCDL::TargetInfo &target,
+ Type type) {
+ return (target.hasFnuzFp8() && isa<Float8E4M3FNUZType>(type)) ||
+ (target.hasOcpFp8() && isa<Float8E4M3FNType>(type));
}
/// Return the `rocdl` intrinsic corresponding to a MFMA operation `mfma`
/// if one exists. This includes checking to ensure the intrinsic is supported
/// on the architecture you are compiling for.
-static std::optional<StringRef> mfmaOpToIntrinsic(MFMAOp mfma,
- Chipset chipset) {
+static std::optional<StringRef>
+mfmaOpToIntrinsic(MFMAOp mfma, const ROCDL::TargetInfo &target) {
uint32_t m = mfma.getM(), n = mfma.getN(), k = mfma.getK(),
b = mfma.getBlocks();
Type sourceElem = getElementTypeOrSelf(mfma.getSourceA().getType());
Type destElem = getElementTypeOrSelf(mfma.getDestC().getType());
if (sourceElem.isF32() && destElem.isF32()) {
- if (mfma.getReducePrecision() && chipset >= kGfx942) {
+ if (mfma.getReducePrecision() &&
+ target.has(llvm::AMDGPU::FEAT_XF32_INSTS)) {
if (m == 32 && n == 32 && k == 4 && b == 1)
return ROCDL::mfma_f32_32x32x4_xf32::getOperationName();
if (m == 16 && n == 16 && k == 8 && b == 1)
@@ -947,7 +907,7 @@ static std::optional<StringRef> mfmaOpToIntrinsic(MFMAOp mfma,
}
if (sourceElem.isF16() && destElem.isF32()) {
- if (chipset >= kGfx950) {
+ if (target.has(llvm::AMDGPU::FEAT_GFX950_INSTS)) {
if (m == 32 && n == 32 && k == 16 && b == 1)
return ROCDL::mfma_f32_32x32x16_f16::getOperationName();
if (m == 16 && n == 16 && k == 32 && b == 1)
@@ -966,13 +926,13 @@ static std::optional<StringRef> mfmaOpToIntrinsic(MFMAOp mfma,
}
if (sourceElem.isBF16() && destElem.isF32()) {
- if (chipset >= kGfx950) {
+ if (target.has(llvm::AMDGPU::FEAT_GFX950_INSTS)) {
if (m == 32 && n == 32 && k == 16 && b == 1)
return ROCDL::mfma_f32_32x32x16_bf16::getOperationName();
if (m == 16 && n == 16 && k == 32 && b == 1)
return ROCDL::mfma_f32_16x16x32_bf16::getOperationName();
}
- if (chipset >= kGfx90a) {
+ if (target.has(llvm::AMDGPU::FEAT_GFX90A_INSTS)) {
if (m == 32 && n == 32 && k == 4 && b == 2)
return ROCDL::mfma_f32_32x32x4bf16_1k::getOperationName();
if (m == 16 && n == 16 && k == 4 && b == 4)
@@ -997,7 +957,7 @@ static std::optional<StringRef> mfmaOpToIntrinsic(MFMAOp mfma,
}
if (sourceElem.isInteger(8) && destElem.isInteger(32)) {
- if (chipset >= kGfx950) {
+ if (target.has(llvm::AMDGPU::FEAT_GFX950_INSTS)) {
if (m == 32 && n == 32 && k == 32 && b == 1)
return ROCDL::mfma_i32_32x32x32_i8::getOperationName();
if (m == 16 && n == 16 && k == 64 && b == 1)
@@ -1013,51 +973,54 @@ static std::optional<StringRef> mfmaOpToIntrinsic(MFMAOp mfma,
return ROCDL::mfma_i32_32x32x8i8::getOperationName();
if (m == 16 && n == 16 && k == 16 && b == 1)
return ROCDL::mfma_i32_16x16x16i8::getOperationName();
- if (m == 32 && n == 32 && k == 16 && b == 1 && chipset >= kGfx942)
+ if (m == 32 && n == 32 && k == 16 && b == 1 &&
+ target.has(llvm::AMDGPU::FEAT_GFX940_INSTS))
return ROCDL::mfma_i32_32x32x16_i8::getOperationName();
- if (m == 16 && n == 16 && k == 32 && b == 1 && chipset >= kGfx942)
+ if (m == 16 && n == 16 && k == 32 && b == 1 &&
+ target.has(llvm::AMDGPU::FEAT_GFX940_INSTS))
return ROCDL::mfma_i32_16x16x32_i8::getOperationName();
}
- if (sourceElem.isF64() && destElem.isF64() && chipset >= kGfx90a) {
+ if (sourceElem.isF64() && destElem.isF64() &&
+ target.has(llvm::AMDGPU::FEAT_GFX90A_INSTS)) {
if (m == 16 && n == 16 && k == 4 && b == 1)
return ROCDL::mfma_f64_16x16x4f64::getOperationName();
if (m == 4 && n == 4 && k == 4 && b == 4)
return ROCDL::mfma_f64_4x4x4f64::getOperationName();
}
- if (destElem.isF32() && typeIsExpectedBf8ForChipset(chipset, sourceElem)) {
+ if (destElem.isF32() && typeIsExpectedBf8ForTarget(target, sourceElem)) {
// Known to be correct because there are no scalar f8 instructions and
// because a length mismatch will have been caught by the verifier.
Type sourceBElem =
cast<VectorType>(mfma.getSourceB().getType()).getElementType();
if (m == 16 && n == 16 && k == 32 && b == 1) {
- if (typeIsExpectedBf8ForChipset(chipset, sourceBElem))
+ if (typeIsExpectedBf8ForTarget(target, sourceBElem))
return ROCDL::mfma_f32_16x16x32_bf8_bf8::getOperationName();
- if (typeIsExpectedFp8ForChipset(chipset, sourceBElem))
+ if (typeIsExpectedFp8ForTarget(target, sourceBElem))
return ROCDL::mfma_f32_16x16x32_bf8_fp8::getOperationName();
}
if (m == 32 && n == 32 && k == 16 && b == 1) {
- if (typeIsExpectedBf8ForChipset(chipset, sourceBElem))
+ if (typeIsExpectedBf8ForTarget(target, sourceBElem))
return ROCDL::mfma_f32_32x32x16_bf8_bf8::getOperationName();
- if (typeIsExpectedFp8ForChipset(chipset, sourceBElem))
+ if (typeIsExpectedFp8ForTarget(target, sourceBElem))
return ROCDL::mfma_f32_32x32x16_bf8_fp8::getOperationName();
}
}
- if (destElem.isF32() && typeIsExpectedFp8ForChipset(chipset, sourceElem)) {
+ if (destElem.isF32() && typeIsExpectedFp8ForTarget(target, sourceElem)) {
Type sourceBElem =
cast<VectorType>(mfma.getSourceB().getType()).getElementType();
if (m == 16 && n == 16 && k == 32 && b == 1) {
- if (typeIsExpectedBf8ForChipset(chipset, sourceBElem))
+ if (typeIsExpectedBf8ForTarget(target, sourceBElem))
return ROCDL::mfma_f32_16x16x32_fp8_bf8::getOperationName();
- if (typeIsExpectedFp8ForChipset(chipset, sourceBElem))
+ if (typeIsExpectedFp8ForTarget(target, sourceBElem))
return ROCDL::mfma_f32_16x16x32_fp8_fp8::getOperationName();
}
if (m == 32 && n == 32 && k == 16 && b == 1) {
- if (typeIsExpectedBf8ForChipset(chipset, sourceBElem))
+ if (typeIsExpectedBf8ForTarget(target, sourceBElem))
return ROCDL::mfma_f32_32x32x16_fp8_bf8::getOperationName();
- if (typeIsExpectedFp8ForChipset(chipset, sourceBElem))
+ if (typeIsExpectedFp8ForTarget(target, sourceBElem))
return ROCDL::mfma_f32_32x32x16_fp8_fp8::getOperationName();
}
}
@@ -1089,12 +1052,13 @@ using ScaledMFMAIntrinsic =
static std::optional<ScaledMFMAIntrinsic>
mfmaOpToScaledIntrinsic(Type aType, Type bType, Type destType, uint32_t m,
- uint32_t n, uint32_t k, uint32_t b, Chipset chipset) {
+ uint32_t n, uint32_t k, uint32_t b,
+ const ROCDL::TargetInfo &target) {
aType = getElementTypeOrSelf(aType);
bType = getElementTypeOrSelf(bType);
destType = getElementTypeOrSelf(destType);
- if (chipset < kGfx950)
+ if (!target.has(llvm::AMDGPU::FEAT_GFX950_INSTS))
return std::nullopt;
if (!isa<Float32Type>(destType))
return std::nullopt;
@@ -1118,19 +1082,19 @@ mfmaOpToScaledIntrinsic(Type aType, Type bType, Type destType, uint32_t m,
}
static std::optional<ScaledMFMAIntrinsic>
-mfmaOpToScaledIntrinsic(MFMAOp mfma, Chipset chipset) {
+mfmaOpToScaledIntrinsic(MFMAOp mfma, const ROCDL::TargetInfo &target) {
return mfmaOpToScaledIntrinsic(
mfma.getSourceA().getType(), mfma.getSourceB().getType(),
mfma.getDestC().getType(), mfma.getM(), mfma.getN(), mfma.getK(),
- mfma.getBlocks(), chipset);
+ mfma.getBlocks(), target);
}
static std::optional<ScaledMFMAIntrinsic>
-mfmaOpToScaledIntrinsic(ScaledMFMAOp smfma, Chipset chipset) {
+mfmaOpToScaledIntrinsic(ScaledMFMAOp smfma, const ROCDL::TargetInfo &target) {
return mfmaOpToScaledIntrinsic(smfma.getSourceA().getType(),
smfma.getSourceB().getType(),
smfma.getDestC().getType(), smfma.getM(),
- smfma.getN(), smfma.getK(), 1u, chipset);
+ smfma.getN(), smfma.getK(), 1u, target);
}
/// Returns the `rocdl` intrinsic corresponding to a WMMA operation `wmma`
@@ -1285,11 +1249,11 @@ static std::optional<StringRef> wmmaOpToIntrinsicGfx1250(Type elemSourceType,
/// Returns the `rocdl` intrinsic corresponding to a SparseMFMA (smfmac)
/// operation if one exists. This includes checking to ensure the intrinsic is
/// supported on the architecture you are compiling for.
-static std::optional<StringRef> smfmacOpToIntrinsic(SparseMFMAOp op,
- Chipset chipset) {
- bool isGfx950 = chipset >= kGfx950;
- auto isFp8 = [&](Type t) { return typeIsExpectedFp8ForChipset(chipset, t); };
- auto isBf8 = [&](Type t) { return typeIsExpectedBf8ForChipset(chipset, t); };
+static std::optional<StringRef>
+smfmacOpToIntrinsic(SparseMFMAOp op, const ROCDL::TargetInfo &target) {
+ bool isGfx950 = target.has(llvm::AMDGPU::FEAT_GFX950_INSTS);
+ auto isFp8 = [&](Type t) { return typeIsExpectedFp8ForTarget(target, t); };
+ auto isBf8 = [&](Type t) { return typeIsExpectedBf8ForTarget(target, t); };
uint32_t m = op.getM(), n = op.getN(), k = op.getK();
Type sourceAElem = getElementTypeOrSelf(op.getSourceA().getType());
@@ -1384,8 +1348,8 @@ static std::optional<StringRef> smfmacOpToIntrinsic(SparseMFMAOp op,
/// Returns the `rocdl` intrinsic corresponding to a WMMA operation `wmma`
/// if one exists. This includes checking to ensure the intrinsic is supported
/// on the architecture you are compiling for.
-static std::optional<StringRef> wmmaOpToIntrinsic(WMMAOp wmma,
- Chipset chipset) {
+static std::optional<StringRef>
+wmmaOpToIntrinsic(WMMAOp wmma, const ROCDL::TargetInfo &target) {
auto sourceVectorType = cast<VectorType>(wmma.getSourceA().getType());
auto sourceBVectorType = cast<VectorType>(wmma.getSourceB().getType());
auto destVectorType = cast<VectorType>(wmma.getDestC().getType());
@@ -1394,8 +1358,9 @@ static std::optional<StringRef> wmmaOpToIntrinsic(WMMAOp wmma,
Type elemDestType = destVectorType.getElementType();
const uint32_t k = wmma.getK();
- const bool isRDNA3 = chipset.majorVersion == 11;
- const bool isRDNA4 = chipset.majorVersion == 12 && chipset.minorVersion == 0;
+ const bool isRDNA3 = target.isGeneration(11);
+ const bool isRDNA4 =
+ target.isGeneration(12) && !target.has(llvm::AMDGPU::FEAT_GFX1250_INSTS);
// Handle RDNA3 and RDNA4.
if (isRDNA3 || isRDNA4)
@@ -1403,7 +1368,7 @@ static std::optional<StringRef> wmmaOpToIntrinsic(WMMAOp wmma,
k, isRDNA3);
// Handle gfx1250.
- if (chipset == kGfx1250)
+ if (target.has(llvm::AMDGPU::FEAT_GFX1250_INSTS))
return wmmaOpToIntrinsicGfx1250(elemSourceType, elemBSourceType,
elemDestType, k);
@@ -1421,7 +1386,7 @@ struct SparseWMMAOpInfo {
};
static std::optional<SparseWMMAOpInfo>
-sparseWMMAOpToIntrinsic(SparseWMMAOp swmmac, Chipset chipset) {
+sparseWMMAOpToIntrinsic(SparseWMMAOp swmmac, const ROCDL::TargetInfo &target) {
Type sourceAElem = getElementTypeOrSelf(swmmac.getSourceA().getType());
Type sourceBElem = getElementTypeOrSelf(swmmac.getSourceB().getType());
Type destElem = getElementTypeOrSelf(swmmac.getDestC().getType());
@@ -1431,7 +1396,8 @@ sparseWMMAOpToIntrinsic(SparseWMMAOp swmmac, Chipset chipset) {
if ((m != 16) || (n != 16))
return std::nullopt;
- const bool isRDNA4 = chipset.majorVersion == 12 && chipset.minorVersion == 0;
+ const bool isRDNA4 =
+ target.isGeneration(12) && !target.has(llvm::AMDGPU::FEAT_GFX1250_INSTS);
if (isRDNA4) {
if (k == 32) {
if (destElem.isF32() && sourceAElem.isF16() && sourceBElem.isF16())
@@ -1489,7 +1455,7 @@ sparseWMMAOpToIntrinsic(SparseWMMAOp swmmac, Chipset chipset) {
}
}
- const bool isGFX1250 = chipset == kGfx1250;
+ const bool isGFX1250 = target.has(llvm::AMDGPU::FEAT_GFX1250_INSTS);
const bool isWavesize64 = swmmac.getWave64();
if (isGFX1250 && !isWavesize64) {
if (k == 64) {
@@ -1567,10 +1533,11 @@ sparseWMMAOpToIntrinsic(SparseWMMAOp swmmac, Chipset chipset) {
namespace {
struct MFMAOpLowering : public ConvertOpToLLVMPattern<MFMAOp> {
- MFMAOpLowering(const LLVMTypeConverter &converter, Chipset chipset)
- : ConvertOpToLLVMPattern<MFMAOp>(converter), chipset(chipset) {}
+ MFMAOpLowering(const LLVMTypeConverter &converter,
+ const ROCDL::TargetInfo &target)
+ : ConvertOpToLLVMPattern<MFMAOp>(converter), target(target) {}
- Chipset chipset;
+ ROCDL::TargetInfo target;
LogicalResult
matchAndRewrite(MFMAOp op, MFMAOpAdaptor adaptor,
@@ -1583,18 +1550,18 @@ struct MFMAOpLowering : public ConvertOpToLLVMPattern<MFMAOp> {
if (outVecType.getElementType().isBF16())
intrinsicOutType = outVecType.clone(rewriter.getI16Type());
- if (chipset.majorVersion != 9 || chipset < kGfx908)
+ if (!target.has(llvm::AMDGPU::FEAT_MAI_INSTS))
return op->emitOpError("MFMA only supported on gfx908+");
uint32_t getBlgpField = static_cast<uint32_t>(op.getBlgp());
if (op.getNegateA() || op.getNegateB() || op.getNegateC()) {
- if (chipset < kGfx942)
+ if (!target.has(llvm::AMDGPU::FEAT_GFX940_INSTS))
return op.emitOpError("negation unsupported on older than gfx942");
getBlgpField |=
op.getNegateA() | (op.getNegateB() << 1) | (op.getNegateC() << 2);
}
- std::optional<StringRef> maybeIntrinsic = mfmaOpToIntrinsic(op, chipset);
+ std::optional<StringRef> maybeIntrinsic = mfmaOpToIntrinsic(op, target);
std::optional<ScaledMFMAIntrinsic> maybeScaledIntrinsic =
- mfmaOpToScaledIntrinsic(op, chipset);
+ mfmaOpToScaledIntrinsic(op, target);
if (!maybeIntrinsic.has_value() && !maybeScaledIntrinsic.has_value())
return op.emitOpError("no intrinsic matching MFMA size on given chipset");
@@ -1612,7 +1579,7 @@ struct MFMAOpLowering : public ConvertOpToLLVMPattern<MFMAOp> {
// Determine if we can use bf16 in the intrinsic. Newer MFMAs in gfx950+
// allows bf16 as the input. For reference check IntrinsicsAMDGPU.td file.
bool allowBf16 = [&]() {
- if (chipset < kGfx950)
+ if (!target.has(llvm::AMDGPU::FEAT_GFX950_INSTS))
return false;
if (isScaled)
return true;
@@ -1660,10 +1627,11 @@ struct MFMAOpLowering : public ConvertOpToLLVMPattern<MFMAOp> {
};
struct ScaledMFMAOpLowering : public ConvertOpToLLVMPattern<ScaledMFMAOp> {
- ScaledMFMAOpLowering(const LLVMTypeConverter &converter, Chipset chipset)
- : ConvertOpToLLVMPattern(converter), chipset(chipset) {}
+ ScaledMFMAOpLowering(const LLVMTypeConverter &converter,
+ const ROCDL::TargetInfo &target)
+ : ConvertOpToLLVMPattern(converter), target(target) {}
- Chipset chipset;
+ ROCDL::TargetInfo target;
LogicalResult
matchAndRewrite(ScaledMFMAOp op, ScaledMFMAOpAdaptor adaptor,
@@ -1671,10 +1639,10 @@ struct ScaledMFMAOpLowering : public ConvertOpToLLVMPattern<ScaledMFMAOp> {
Location loc = op.getLoc();
Type intrinsicOutType = typeConverter->convertType(op.getDestD().getType());
- if (chipset.majorVersion != 9 || chipset < kGfx950)
+ if (!target.has(llvm::AMDGPU::FEAT_GFX950_INSTS))
return op->emitOpError("scaled MFMA only supported on gfx908+");
std::optional<ScaledMFMAIntrinsic> maybeScaledIntrinsic =
- mfmaOpToScaledIntrinsic(op, chipset);
+ mfmaOpToScaledIntrinsic(op, target);
if (!maybeScaledIntrinsic.has_value())
return op.emitOpError(
"no intrinsic matching scaled MFMA size on given chipset");
@@ -1706,10 +1674,11 @@ struct ScaledMFMAOpLowering : public ConvertOpToLLVMPattern<ScaledMFMAOp> {
};
struct SparseMFMAOpLowering : public ConvertOpToLLVMPattern<SparseMFMAOp> {
- SparseMFMAOpLowering(const LLVMTypeConverter &converter, Chipset chipset)
- : ConvertOpToLLVMPattern<SparseMFMAOp>(converter), chipset(chipset) {}
+ SparseMFMAOpLowering(const LLVMTypeConverter &converter,
+ const ROCDL::TargetInfo &target)
+ : ConvertOpToLLVMPattern<SparseMFMAOp>(converter), target(target) {}
- Chipset chipset;
+ ROCDL::TargetInfo target;
LogicalResult
matchAndRewrite(SparseMFMAOp op, SparseMFMAOpAdaptor adaptor,
@@ -1721,10 +1690,10 @@ struct SparseMFMAOpLowering : public ConvertOpToLLVMPattern<SparseMFMAOp> {
return rewriter.notifyMatchFailure(op, "type conversion failed");
// smfmac is supported on gfx942 and gfx950.
- if (chipset.majorVersion != 9 || chipset < kGfx942)
+ if (!target.has(llvm::AMDGPU::FEAT_GFX940_INSTS))
return op->emitOpError("sparse MFMA (smfmac) only supported on gfx942+");
- std::optional<StringRef> maybeIntrinsic = smfmacOpToIntrinsic(op, chipset);
+ std::optional<StringRef> maybeIntrinsic = smfmacOpToIntrinsic(op, target);
if (!maybeIntrinsic.has_value())
return op.emitOpError(
"no intrinsic matching sparse MFMA on the given chipset");
@@ -1733,7 +1702,8 @@ struct SparseMFMAOpLowering : public ConvertOpToLLVMPattern<SparseMFMAOp> {
ROCDL::smfmac_f32_16x16x32_bf16::getOperationName() ||
*maybeIntrinsic ==
ROCDL::smfmac_f32_32x32x16_bf16::getOperationName());
- bool isGfx950 = (chipset >= kGfx950) && !isGfx942BF16;
+ bool isGfx950 =
+ (target.has(llvm::AMDGPU::FEAT_GFX950_INSTS)) && !isGfx942BF16;
Value a = convertPackedVectorOperand(rewriter, loc, adaptor.getSourceA(),
isGfx950);
@@ -1761,10 +1731,11 @@ struct SparseMFMAOpLowering : public ConvertOpToLLVMPattern<SparseMFMAOp> {
};
struct WMMAOpLowering : public ConvertOpToLLVMPattern<WMMAOp> {
- WMMAOpLowering(const LLVMTypeConverter &converter, Chipset chipset)
- : ConvertOpToLLVMPattern<WMMAOp>(converter), chipset(chipset) {}
+ WMMAOpLowering(const LLVMTypeConverter &converter,
+ const ROCDL::TargetInfo &target)
+ : ConvertOpToLLVMPattern<WMMAOp>(converter), target(target) {}
- Chipset chipset;
+ ROCDL::TargetInfo target;
LogicalResult
matchAndRewrite(WMMAOp op, WMMAOpAdaptor adaptor,
@@ -1775,10 +1746,10 @@ struct WMMAOpLowering : public ConvertOpToLLVMPattern<WMMAOp> {
if (!outType)
return rewriter.notifyMatchFailure(op, "type conversion failed");
- if (chipset.majorVersion != 11 && chipset.majorVersion != 12)
+ if (!target.isGeneration(11) && !target.isGeneration(12))
return op->emitOpError("WMMA only supported on gfx11 and gfx12");
- bool isGFX1250 = chipset >= kGfx1250;
+ bool isGFX1250 = target.has(llvm::AMDGPU::FEAT_GFX1250_INSTS);
// The WMMA operations represent vectors of bf16s as vectors of i16s
// (except on gfx1250), so we need to bitcast bfloats to i16 and then
@@ -1806,12 +1777,13 @@ struct WMMAOpLowering : public ConvertOpToLLVMPattern<WMMAOp> {
destC = LLVM::BitcastOp::create(
rewriter, loc, destCType.clone(rewriter.getI16Type()), destC);
- std::optional<StringRef> maybeIntrinsic = wmmaOpToIntrinsic(op, chipset);
+ std::optional<StringRef> maybeIntrinsic = wmmaOpToIntrinsic(op, target);
if (!maybeIntrinsic.has_value())
return op.emitOpError("no intrinsic matching WMMA on the given chipset");
- if (chipset.majorVersion >= 12 && op.getSubwordOffset() != 0)
+ if (target.has(llvm::AMDGPU::FEAT_GFX12_INSTS) &&
+ op.getSubwordOffset() != 0)
return op.emitOpError("subwordOffset not supported on gfx12+");
SmallVector<Value, 4> operands;
@@ -1850,7 +1822,7 @@ enum class DotFamily {
};
static std::optional<std::pair<StringRef, DotFamily>>
-dotOpToIntrinsic(DotOp op, Chipset chipset) {
+dotOpToIntrinsic(DotOp op, const ROCDL::TargetInfo &target) {
Type aElem = cast<VectorType>(op.getSourceA().getType()).getElementType();
Type bElem = cast<VectorType>(op.getSourceB().getType()).getElementType();
Type dest = op.getDestC().getType();
@@ -1859,18 +1831,18 @@ dotOpToIntrinsic(DotOp op, Chipset chipset) {
// f16 x f16 -> f32 / f16.
if (aElem.isF16() && bElem.isF16()) {
- if (dest.isF32() && hasDot10Insts(chipset))
+ if (dest.isF32() && target.has(llvm::AMDGPU::FEAT_DOT10_INSTS))
return {{ROCDL::fdot2::getOperationName(), DotFamily::Clamp}};
- if (dest.isF16() && hasDot9Insts(chipset))
+ if (dest.isF16() && target.has(llvm::AMDGPU::FEAT_DOT9_INSTS))
return {{ROCDL::fdot2_f16_f16::getOperationName(), DotFamily::NoClamp}};
return std::nullopt;
}
// bf16 x bf16 -> f32 / bf16.
if (aElem.isBF16() && bElem.isBF16()) {
- if (dest.isF32() && hasDot12Insts(chipset))
+ if (dest.isF32() && target.has(llvm::AMDGPU::FEAT_DOT12_INSTS))
return {{ROCDL::fdot2_f32_bf16::getOperationName(), DotFamily::Clamp}};
- if (dest.isBF16() && hasDot9Insts(chipset))
+ if (dest.isBF16() && target.has(llvm::AMDGPU::FEAT_DOT9_INSTS))
return {{ROCDL::fdot2_bf16_bf16::getOperationName(), DotFamily::NoClamp}};
return std::nullopt;
}
@@ -1882,7 +1854,7 @@ dotOpToIntrinsic(DotOp op, Chipset chipset) {
unsigned elemWidth = aElem.getIntOrFloatBitWidth();
if (mixedSign) {
- if (!hasDot8Insts(chipset))
+ if (!target.has(llvm::AMDGPU::FEAT_DOT8_INSTS))
return std::nullopt;
StringRef name;
switch (elemWidth) {
@@ -1902,19 +1874,21 @@ dotOpToIntrinsic(DotOp op, Chipset chipset) {
bool supported = false;
switch (elemWidth) {
case 16:
- supported = hasDot2Insts(chipset);
+ supported = target.has(llvm::AMDGPU::FEAT_DOT2_INSTS);
name = uA ? ROCDL::udot2::getOperationName()
: ROCDL::sdot2::getOperationName();
break;
case 8:
- supported = uA ? hasDot7Insts(chipset)
- : hasDot1Insts(chipset) || hasDot8Insts(chipset);
+ supported = uA ? target.has(llvm::AMDGPU::FEAT_DOT7_INSTS)
+ : target.has(llvm::AMDGPU::FEAT_DOT1_INSTS) ||
+ target.has(llvm::AMDGPU::FEAT_DOT8_INSTS);
name = uA ? ROCDL::udot4::getOperationName()
: ROCDL::sdot4::getOperationName();
break;
case 4:
- supported = uA ? hasDot7Insts(chipset)
- : hasDot1Insts(chipset) || hasDot8Insts(chipset);
+ supported = uA ? target.has(llvm::AMDGPU::FEAT_DOT7_INSTS)
+ : target.has(llvm::AMDGPU::FEAT_DOT1_INSTS) ||
+ target.has(llvm::AMDGPU::FEAT_DOT8_INSTS);
name = uA ? ROCDL::udot8::getOperationName()
: ROCDL::sdot8::getOperationName();
break;
@@ -1932,7 +1906,7 @@ dotOpToIntrinsic(DotOp op, Chipset chipset) {
bool bIsFp8 = isa<Float8E4M3FNType>(bElem);
bool bIsBf8 = isa<Float8E5M2Type>(bElem);
if ((aIsFp8 || aIsBf8) && (bIsFp8 || bIsBf8) && dest.isF32()) {
- if (!hasDot11Insts(chipset))
+ if (!target.has(llvm::AMDGPU::FEAT_DOT11_INSTS))
return std::nullopt;
StringRef name;
if (aIsFp8 && bIsFp8)
@@ -1950,10 +1924,11 @@ dotOpToIntrinsic(DotOp op, Chipset chipset) {
}
struct DotOpLowering : public ConvertOpToLLVMPattern<DotOp> {
- DotOpLowering(const LLVMTypeConverter &converter, Chipset chipset)
- : ConvertOpToLLVMPattern<DotOp>(converter), chipset(chipset) {}
+ DotOpLowering(const LLVMTypeConverter &converter,
+ const ROCDL::TargetInfo &target)
+ : ConvertOpToLLVMPattern<DotOp>(converter), target(target) {}
- Chipset chipset;
+ ROCDL::TargetInfo target;
LogicalResult
matchAndRewrite(DotOp op, DotOpAdaptor adaptor,
@@ -1961,7 +1936,7 @@ struct DotOpLowering : public ConvertOpToLLVMPattern<DotOp> {
Location loc = op.getLoc();
std::optional<std::pair<StringRef, DotFamily>> maybeIntrinsic =
- dotOpToIntrinsic(op, chipset);
+ dotOpToIntrinsic(op, target);
if (!maybeIntrinsic)
return op.emitOpError("no intrinsic matching dot on the given chipset: ")
<< op.getSourceA().getType() << " * " << op.getSourceB().getType()
@@ -1998,10 +1973,11 @@ struct DotOpLowering : public ConvertOpToLLVMPattern<DotOp> {
};
struct SparseWMMAOpLowering : public ConvertOpToLLVMPattern<SparseWMMAOp> {
- SparseWMMAOpLowering(const LLVMTypeConverter &converter, Chipset chipset)
- : ConvertOpToLLVMPattern<SparseWMMAOp>(converter), chipset(chipset) {}
+ SparseWMMAOpLowering(const LLVMTypeConverter &converter,
+ const ROCDL::TargetInfo &target)
+ : ConvertOpToLLVMPattern<SparseWMMAOp>(converter), target(target) {}
- Chipset chipset;
+ ROCDL::TargetInfo target;
LogicalResult
matchAndRewrite(SparseWMMAOp op, SparseWMMAOpAdaptor adaptor,
@@ -2013,7 +1989,7 @@ struct SparseWMMAOpLowering : public ConvertOpToLLVMPattern<SparseWMMAOp> {
return rewriter.notifyMatchFailure(op, "type conversion failed");
std::optional<SparseWMMAOpInfo> maybeIntrinsic =
- sparseWMMAOpToIntrinsic(op, chipset);
+ sparseWMMAOpToIntrinsic(op, target);
if (!maybeIntrinsic.has_value())
return op.emitOpError(
@@ -2045,8 +2021,7 @@ struct SparseWMMAOpLowering : public ConvertOpToLLVMPattern<SparseWMMAOp> {
if (intrinsic.useClamp && op.getClampAttr())
attrs.push_back({"clamp", op.getClampAttr()});
- const bool isGFX1250orHigher =
- chipset.majorVersion == 12 && chipset.minorVersion >= 5;
+ const bool isGFX1250orHigher = target.has(llvm::AMDGPU::FEAT_GFX1250_INSTS);
Value a = convertPackedVectorOperand(rewriter, loc, adaptor.getSourceA(),
isGFX1250orHigher);
Value b = convertPackedVectorOperand(rewriter, loc, adaptor.getSourceB(),
@@ -2079,10 +2054,11 @@ struct SparseWMMAOpLowering : public ConvertOpToLLVMPattern<SparseWMMAOp> {
};
struct ScaledWMMAOpLowering : public ConvertOpToLLVMPattern<ScaledWMMAOp> {
- ScaledWMMAOpLowering(const LLVMTypeConverter &converter, Chipset chipset)
- : ConvertOpToLLVMPattern<ScaledWMMAOp>(converter), chipset(chipset) {}
+ ScaledWMMAOpLowering(const LLVMTypeConverter &converter,
+ const ROCDL::TargetInfo &target)
+ : ConvertOpToLLVMPattern<ScaledWMMAOp>(converter), target(target) {}
- Chipset chipset;
+ ROCDL::TargetInfo target;
LogicalResult
matchAndRewrite(ScaledWMMAOp op, ScaledWMMAOpAdaptor adaptor,
@@ -2093,7 +2069,7 @@ struct ScaledWMMAOpLowering : public ConvertOpToLLVMPattern<ScaledWMMAOp> {
if (!outType)
return rewriter.notifyMatchFailure(op, "type conversion failed");
- if (chipset < kGfx1250)
+ if (!target.has(llvm::AMDGPU::FEAT_GFX1250_INSTS))
return op->emitOpError("WMMA scale only supported on gfx1250+");
int64_t m = op.getM();
@@ -2199,15 +2175,17 @@ struct ScaledWMMAOpLowering : public ConvertOpToLLVMPattern<ScaledWMMAOp> {
struct TransposeLoadOpLowering
: public ConvertOpToLLVMPattern<TransposeLoadOp> {
- TransposeLoadOpLowering(const LLVMTypeConverter &converter, Chipset chipset)
- : ConvertOpToLLVMPattern<TransposeLoadOp>(converter), chipset(chipset) {}
+ TransposeLoadOpLowering(const LLVMTypeConverter &converter,
+ const ROCDL::TargetInfo &target)
+ : ConvertOpToLLVMPattern<TransposeLoadOp>(converter), target(target) {}
- Chipset chipset;
+ ROCDL::TargetInfo target;
LogicalResult
matchAndRewrite(TransposeLoadOp op, TransposeLoadOpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
- if (chipset != kGfx950 && chipset < kGfx1250)
+ if (!target.has(llvm::AMDGPU::FEAT_GFX950_INSTS) &&
+ !target.has(llvm::AMDGPU::FEAT_GFX1250_INSTS))
return op.emitOpError(
"transpose_load is only supported on gfx950 and gfx1250+");
@@ -2248,7 +2226,7 @@ struct TransposeLoadOpLowering
};
Value intrinsic;
- if (chipset >= kGfx1250) {
+ if (target.has(llvm::AMDGPU::FEAT_GFX1250_INSTS)) {
switch (elementTypeSize) {
case 4: {
if (numElements != 16)
@@ -2348,17 +2326,17 @@ struct TransposeLoadOpLowering
struct GlobalTransposeLoadOpLowering
: public ConvertOpToLLVMPattern<GlobalTransposeLoadOp> {
GlobalTransposeLoadOpLowering(const LLVMTypeConverter &converter,
- Chipset chipset)
+ const ROCDL::TargetInfo &target)
: ConvertOpToLLVMPattern<GlobalTransposeLoadOp>(converter),
- chipset(chipset) {}
+ target(target) {}
- Chipset chipset;
+ ROCDL::TargetInfo target;
LogicalResult
matchAndRewrite(GlobalTransposeLoadOp op,
GlobalTransposeLoadOpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
- if (chipset < kGfx1200)
+ if (!target.has(llvm::AMDGPU::FEAT_GFX12_INSTS))
return op.emitOpError(
"global_transpose_load is only supported on gfx1200+");
@@ -2386,7 +2364,7 @@ struct GlobalTransposeLoadOpLowering
switch (elementTypeSize) {
case 4: {
assert(numElements == 16);
- if (chipset < kGfx1250)
+ if (!target.has(llvm::AMDGPU::FEAT_GFX1250_INSTS))
return op.emitOpError("4-bit global_transpose_load requires gfx1250+");
auto rocdlOp = ROCDL::GlobalLoadTr4_B64::create(
rewriter, loc, rocdlResultType, srcPtr, ArrayAttr{}, ArrayAttr{},
@@ -2396,7 +2374,7 @@ struct GlobalTransposeLoadOpLowering
}
case 6: {
assert(numElements == 16);
- if (chipset < kGfx1250)
+ if (!target.has(llvm::AMDGPU::FEAT_GFX1250_INSTS))
return op.emitOpError("6-bit global_transpose_load requires gfx1250+");
auto rocdlOp = ROCDL::GlobalLoadTr6_B96::create(
rewriter, loc, rocdlResultType, srcPtr, ArrayAttr{}, ArrayAttr{},
@@ -2427,15 +2405,16 @@ struct GlobalTransposeLoadOpLowering
};
struct GatherToLDSOpLowering : public ConvertOpToLLVMPattern<GatherToLDSOp> {
- GatherToLDSOpLowering(const LLVMTypeConverter &converter, Chipset chipset)
- : ConvertOpToLLVMPattern<GatherToLDSOp>(converter), chipset(chipset) {}
+ GatherToLDSOpLowering(const LLVMTypeConverter &converter,
+ const ROCDL::TargetInfo &target)
+ : ConvertOpToLLVMPattern<GatherToLDSOp>(converter), target(target) {}
- Chipset chipset;
+ ROCDL::TargetInfo target;
LogicalResult
matchAndRewrite(GatherToLDSOp op, GatherToLDSOpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
- if (chipset.majorVersion < 9 || chipset.majorVersion > 10)
+ if (!target.has(llvm::AMDGPU::FEAT_VMEM_TO_LDS_LOAD_INSTS))
return op.emitOpError("pre-gfx9 and post-gfx10 not supported");
Location loc = op.getLoc();
@@ -2460,7 +2439,8 @@ struct GatherToLDSOpLowering : public ConvertOpToLLVMPattern<GatherToLDSOp> {
if (!llvm::is_contained({1, 2, 4, 12, 16}, loadWidth))
return op.emitOpError("chipset unsupported element size");
- if (chipset != kGfx950 && llvm::is_contained({12, 16}, loadWidth))
+ if (!target.has(llvm::AMDGPU::FEAT_GFX950_INSTS) &&
+ llvm::is_contained({12, 16}, loadWidth))
return op.emitOpError("Gather to LDS instructions with 12-byte and "
"16-byte load widths are only supported on gfx950");
@@ -2492,17 +2472,17 @@ struct GatherToLDSOpLowering : public ConvertOpToLLVMPattern<GatherToLDSOp> {
struct GlobalLoadAsyncToLDSOpLowering
: public ConvertOpToLLVMPattern<GlobalLoadAsyncToLDSOp> {
GlobalLoadAsyncToLDSOpLowering(const LLVMTypeConverter &converter,
- Chipset chipset)
+ const ROCDL::TargetInfo &target)
: ConvertOpToLLVMPattern<GlobalLoadAsyncToLDSOp>(converter),
- chipset(chipset) {}
+ target(target) {}
- Chipset chipset;
+ ROCDL::TargetInfo target;
LogicalResult
matchAndRewrite(GlobalLoadAsyncToLDSOp op,
GlobalLoadAsyncToLDSOpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
- if (chipset < kGfx1250)
+ if (!target.has(llvm::AMDGPU::FEAT_GFX1250_INSTS))
return op.emitOpError(
"global_load_async_to_lds is only supported on gfx1250+");
@@ -2569,10 +2549,11 @@ struct GlobalLoadAsyncToLDSOpLowering
namespace {
struct ExtPackedFp8OpLowering final
: public ConvertOpToLLVMPattern<ExtPackedFp8Op> {
- ExtPackedFp8OpLowering(const LLVMTypeConverter &converter, Chipset chipset)
+ ExtPackedFp8OpLowering(const LLVMTypeConverter &converter,
+ const ROCDL::TargetInfo &target)
: ConvertOpToLLVMPattern<amdgpu::ExtPackedFp8Op>(converter),
- chipset(chipset) {}
- Chipset chipset;
+ target(target) {}
+ ROCDL::TargetInfo target;
LogicalResult
matchAndRewrite(ExtPackedFp8Op op, ExtPackedFp8OpAdaptor adaptor,
@@ -2582,10 +2563,10 @@ struct ExtPackedFp8OpLowering final
struct ScaledExtPackedMatrixOpLowering final
: public ConvertOpToLLVMPattern<ScaledExtPackedMatrixOp> {
ScaledExtPackedMatrixOpLowering(const LLVMTypeConverter &converter,
- Chipset chipset)
+ const ROCDL::TargetInfo &target)
: ConvertOpToLLVMPattern<amdgpu::ScaledExtPackedMatrixOp>(converter),
- chipset(chipset) {}
- Chipset chipset;
+ target(target) {}
+ ROCDL::TargetInfo target;
LogicalResult
matchAndRewrite(ScaledExtPackedMatrixOp op,
@@ -2596,10 +2577,10 @@ struct ScaledExtPackedMatrixOpLowering final
struct PackedTrunc2xFp8OpLowering final
: public ConvertOpToLLVMPattern<PackedTrunc2xFp8Op> {
PackedTrunc2xFp8OpLowering(const LLVMTypeConverter &converter,
- Chipset chipset)
+ const ROCDL::TargetInfo &target)
: ConvertOpToLLVMPattern<amdgpu::PackedTrunc2xFp8Op>(converter),
- chipset(chipset) {}
- Chipset chipset;
+ target(target) {}
+ ROCDL::TargetInfo target;
LogicalResult
matchAndRewrite(PackedTrunc2xFp8Op op, PackedTrunc2xFp8OpAdaptor adaptor,
@@ -2609,10 +2590,10 @@ struct PackedTrunc2xFp8OpLowering final
struct PackedStochRoundFp8OpLowering final
: public ConvertOpToLLVMPattern<PackedStochRoundFp8Op> {
PackedStochRoundFp8OpLowering(const LLVMTypeConverter &converter,
- Chipset chipset)
+ const ROCDL::TargetInfo &target)
: ConvertOpToLLVMPattern<amdgpu::PackedStochRoundFp8Op>(converter),
- chipset(chipset) {}
- Chipset chipset;
+ target(target) {}
+ ROCDL::TargetInfo target;
LogicalResult
matchAndRewrite(PackedStochRoundFp8Op op,
@@ -2622,10 +2603,11 @@ struct PackedStochRoundFp8OpLowering final
struct ScaledExtPackedOpLowering final
: public ConvertOpToLLVMPattern<ScaledExtPackedOp> {
- ScaledExtPackedOpLowering(const LLVMTypeConverter &converter, Chipset chipset)
+ ScaledExtPackedOpLowering(const LLVMTypeConverter &converter,
+ const ROCDL::TargetInfo &target)
: ConvertOpToLLVMPattern<amdgpu::ScaledExtPackedOp>(converter),
- chipset(chipset) {}
- Chipset chipset;
+ target(target) {}
+ ROCDL::TargetInfo target;
LogicalResult
matchAndRewrite(ScaledExtPackedOp op, ScaledExtPackedOpAdaptor adaptor,
@@ -2635,10 +2617,10 @@ struct ScaledExtPackedOpLowering final
struct PackedScaledTruncOpLowering final
: public ConvertOpToLLVMPattern<PackedScaledTruncOp> {
PackedScaledTruncOpLowering(const LLVMTypeConverter &converter,
- Chipset chipset)
+ const ROCDL::TargetInfo &target)
: ConvertOpToLLVMPattern<amdgpu::PackedScaledTruncOp>(converter),
- chipset(chipset) {}
- Chipset chipset;
+ target(target) {}
+ ROCDL::TargetInfo target;
LogicalResult
matchAndRewrite(PackedScaledTruncOp op, PackedScaledTruncOpAdaptor adaptor,
@@ -2651,7 +2633,7 @@ LogicalResult ExtPackedFp8OpLowering::matchAndRewrite(
ExtPackedFp8Op op, ExtPackedFp8OpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const {
Location loc = op.getLoc();
- if (!(chipset == kGfx942 || hasOcpFp8(chipset)))
+ if (!target.has(llvm::AMDGPU::FEAT_FP8_CONVERSION_INSTS))
return rewriter.notifyMatchFailure(
loc, "Fp8 conversion instructions are not available on target "
"architecture and their emulation is not implemented");
@@ -2682,18 +2664,18 @@ LogicalResult ExtPackedFp8OpLowering::matchAndRewrite(
}
Value i32Source = LLVM::BitcastOp::create(rewriter, loc, i32, source);
if (resultVecType) {
- if (typeIsExpectedBf8ForChipset(chipset, sourceElemType)) {
+ if (typeIsExpectedBf8ForTarget(target, sourceElemType)) {
rewriter.replaceOpWithNewOp<ROCDL::CvtPkF32Bf8Op>(op, f32, i32Source,
op.getIndex());
- } else if (typeIsExpectedFp8ForChipset(chipset, sourceElemType)) {
+ } else if (typeIsExpectedFp8ForTarget(target, sourceElemType)) {
rewriter.replaceOpWithNewOp<ROCDL::CvtPkF32Fp8Op>(op, f32, i32Source,
op.getIndex());
}
} else {
- if (typeIsExpectedBf8ForChipset(chipset, sourceElemType)) {
+ if (typeIsExpectedBf8ForTarget(target, sourceElemType)) {
rewriter.replaceOpWithNewOp<ROCDL::CvtF32Bf8Op>(op, f32, i32Source,
op.getIndex());
- } else if (typeIsExpectedFp8ForChipset(chipset, sourceElemType)) {
+ } else if (typeIsExpectedFp8ForTarget(target, sourceElemType)) {
rewriter.replaceOpWithNewOp<ROCDL::CvtF32Fp8Op>(op, f32, i32Source,
op.getIndex());
}
@@ -2801,7 +2783,7 @@ LogicalResult ScaledExtPackedMatrixOpLowering::matchAndRewrite(
using fp6 = Float6E2M3FNType;
using bf6 = Float6E3M2FNType;
Location loc = op.getLoc();
- if (chipset != kGfx1250) {
+ if (!target.has(llvm::AMDGPU::FEAT_GFX1250_INSTS)) {
return rewriter.notifyMatchFailure(
loc,
"Scaled fp packed conversion instructions are not available on target "
@@ -2872,7 +2854,7 @@ LogicalResult ScaledExtPackedOpLowering::matchAndRewrite(
ScaledExtPackedOp op, ScaledExtPackedOpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const {
Location loc = op.getLoc();
- if (chipset != kGfx950)
+ if (!target.has(llvm::AMDGPU::FEAT_GFX950_INSTS))
return rewriter.notifyMatchFailure(
loc, "Scaled fp conversion instructions are not available on target "
"architecture and their emulation is not implemented");
@@ -2952,7 +2934,7 @@ LogicalResult PackedScaledTruncOpLowering::matchAndRewrite(
PackedScaledTruncOp op, PackedScaledTruncOpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const {
Location loc = op.getLoc();
- if (chipset != kGfx950)
+ if (!target.has(llvm::AMDGPU::FEAT_GFX950_INSTS))
return rewriter.notifyMatchFailure(
loc, "Scaled fp conversion instructions are not available on target "
"architecture and their emulation is not implemented");
@@ -3034,7 +3016,7 @@ LogicalResult PackedTrunc2xFp8OpLowering::matchAndRewrite(
PackedTrunc2xFp8Op op, PackedTrunc2xFp8OpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const {
Location loc = op.getLoc();
- if (!(chipset == kGfx942 || hasOcpFp8(chipset)))
+ if (!target.has(llvm::AMDGPU::FEAT_FP8_CONVERSION_INSTS))
return rewriter.notifyMatchFailure(
loc, "Fp8 conversion instructions are not available on target "
"architecture and their emulation is not implemented");
@@ -3054,10 +3036,10 @@ LogicalResult PackedTrunc2xFp8OpLowering::matchAndRewrite(
existing = LLVM::UndefOp::create(rewriter, loc, i32);
Value result;
- if (typeIsExpectedBf8ForChipset(chipset, resultElemType))
+ if (typeIsExpectedBf8ForTarget(target, resultElemType))
result = ROCDL::CvtPkBf8F32Op::create(rewriter, loc, i32, sourceA, sourceB,
existing, op.getWordIndex());
- else if (typeIsExpectedFp8ForChipset(chipset, resultElemType))
+ else if (typeIsExpectedFp8ForTarget(target, resultElemType))
result = ROCDL::CvtPkFp8F32Op::create(rewriter, loc, i32, sourceA, sourceB,
existing, op.getWordIndex());
else
@@ -3073,7 +3055,7 @@ LogicalResult PackedStochRoundFp8OpLowering::matchAndRewrite(
PackedStochRoundFp8Op op, PackedStochRoundFp8OpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const {
Location loc = op.getLoc();
- if (!(chipset == kGfx942 || hasOcpFp8(chipset)))
+ if (!target.has(llvm::AMDGPU::FEAT_FP8_CONVERSION_INSTS))
return rewriter.notifyMatchFailure(
loc, "Fp8 conversion instructions are not available on target "
"architecture and their emulation is not implemented");
@@ -3091,10 +3073,10 @@ LogicalResult PackedStochRoundFp8OpLowering::matchAndRewrite(
existing = LLVM::UndefOp::create(rewriter, loc, i32);
Value result;
- if (typeIsExpectedBf8ForChipset(chipset, resultElemType))
+ if (typeIsExpectedBf8ForTarget(target, resultElemType))
result = ROCDL::CvtSrBf8F32Op::create(rewriter, loc, i32, source, stoch,
existing, op.getStoreIndex());
- else if (typeIsExpectedFp8ForChipset(chipset, resultElemType))
+ else if (typeIsExpectedFp8ForTarget(target, resultElemType))
result = ROCDL::CvtSrFp8F32Op::create(rewriter, loc, i32, source, stoch,
existing, op.getStoreIndex());
else
@@ -3109,9 +3091,10 @@ LogicalResult PackedStochRoundFp8OpLowering::matchAndRewrite(
// Implement the AMDGPU_DPPLowering class that will convert the amdgpu.dpp
// operation into the corresponding ROCDL instructions.
struct AMDGPUDPPLowering : public ConvertOpToLLVMPattern<DPPOp> {
- AMDGPUDPPLowering(const LLVMTypeConverter &converter, Chipset chipset)
- : ConvertOpToLLVMPattern<DPPOp>(converter), chipset(chipset) {}
- Chipset chipset;
+ AMDGPUDPPLowering(const LLVMTypeConverter &converter,
+ const ROCDL::TargetInfo &target)
+ : ConvertOpToLLVMPattern<DPPOp>(converter), target(target) {}
+ ROCDL::TargetInfo target;
LogicalResult
matchAndRewrite(DPPOp DppOp, DPPOp::Adaptor adaptor,
@@ -3295,20 +3278,25 @@ struct AMDGPUSwizzleBitModeLowering
struct AMDGPUPermlaneLowering : public ConvertOpToLLVMPattern<PermlaneSwapOp> {
using ConvertOpToLLVMPattern::ConvertOpToLLVMPattern;
- AMDGPUPermlaneLowering(const LLVMTypeConverter &converter, Chipset chipset)
- : ConvertOpToLLVMPattern<PermlaneSwapOp>(converter), chipset(chipset) {}
- Chipset chipset;
+ AMDGPUPermlaneLowering(const LLVMTypeConverter &converter,
+ const ROCDL::TargetInfo &target)
+ : ConvertOpToLLVMPattern<PermlaneSwapOp>(converter), target(target) {}
+ ROCDL::TargetInfo target;
LogicalResult
matchAndRewrite(PermlaneSwapOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
- if (chipset < kGfx950)
- return op->emitOpError("permlane_swap is only supported on gfx950+");
+ unsigned rowLength = op.getRowLength();
+ bool supported = rowLength == 16
+ ? target.has(llvm::AMDGPU::FEAT_PERMLANE16_SWAP)
+ : target.has(llvm::AMDGPU::FEAT_PERMLANE32_SWAP);
+ if (!supported)
+ return op->emitOpError("permlane_swap of row length ")
+ << rowLength << " is not supported on " << target.getArchName();
Location loc = op.getLoc();
Type i32 = rewriter.getI32Type();
Value src = adaptor.getSrc();
- unsigned rowLength = op.getRowLength();
bool fi = op.getFetchInactive();
bool boundctrl = op.getBoundCtrl();
@@ -3355,14 +3343,15 @@ struct AMDGPUPermlaneVarLowering
: public ConvertOpToLLVMPattern<PermlaneVarOp> {
using ConvertOpToLLVMPattern::ConvertOpToLLVMPattern;
- AMDGPUPermlaneVarLowering(const LLVMTypeConverter &converter, Chipset chipset)
- : ConvertOpToLLVMPattern<PermlaneVarOp>(converter), chipset(chipset) {}
- Chipset chipset;
+ AMDGPUPermlaneVarLowering(const LLVMTypeConverter &converter,
+ const ROCDL::TargetInfo &target)
+ : ConvertOpToLLVMPattern<PermlaneVarOp>(converter), target(target) {}
+ ROCDL::TargetInfo target;
LogicalResult
matchAndRewrite(PermlaneVarOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
- if (chipset < kGfx1200)
+ if (!target.has(llvm::AMDGPU::FEAT_GFX12_INSTS))
return op->emitOpError("permlane_var is only supported on GFX12+");
Location loc = op.getLoc();
@@ -3412,15 +3401,16 @@ constexpr int32_t kDsBarrierPendingCountMask =
struct DsBarrierInitOpLowering
: public ConvertOpToLLVMPattern<DsBarrierInitOp> {
- Chipset chipset;
+ ROCDL::TargetInfo target;
- DsBarrierInitOpLowering(const LLVMTypeConverter &converter, Chipset chipset)
- : ConvertOpToLLVMPattern<DsBarrierInitOp>(converter), chipset(chipset) {}
+ DsBarrierInitOpLowering(const LLVMTypeConverter &converter,
+ const ROCDL::TargetInfo &target)
+ : ConvertOpToLLVMPattern<DsBarrierInitOp>(converter), target(target) {}
LogicalResult
matchAndRewrite(DsBarrierInitOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
- if (chipset < kGfx1250)
+ if (!target.has(llvm::AMDGPU::FEAT_GFX1250_INSTS))
return op->emitOpError("only supported on gfx1250+");
Location loc = op.getLoc();
@@ -3465,17 +3455,17 @@ struct DsBarrierInitOpLowering
struct DsBarrierPollStateOpLowering
: public ConvertOpToLLVMPattern<DsBarrierPollStateOp> {
- Chipset chipset;
+ ROCDL::TargetInfo target;
DsBarrierPollStateOpLowering(const LLVMTypeConverter &converter,
- Chipset chipset)
+ const ROCDL::TargetInfo &target)
: ConvertOpToLLVMPattern<DsBarrierPollStateOp>(converter),
- chipset(chipset) {}
+ target(target) {}
LogicalResult
matchAndRewrite(DsBarrierPollStateOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
- if (chipset < kGfx1250)
+ if (!target.has(llvm::AMDGPU::FEAT_GFX1250_INSTS))
return op->emitOpError("only supported on gfx1250+");
Location loc = op.getLoc();
@@ -3498,17 +3488,17 @@ struct DsBarrierPollStateOpLowering
struct DsAsyncBarrierArriveOpLowering
: public ConvertOpToLLVMPattern<DsAsyncBarrierArriveOp> {
- Chipset chipset;
+ ROCDL::TargetInfo target;
DsAsyncBarrierArriveOpLowering(const LLVMTypeConverter &converter,
- Chipset chipset)
+ const ROCDL::TargetInfo &target)
: ConvertOpToLLVMPattern<DsAsyncBarrierArriveOp>(converter),
- chipset(chipset) {}
+ target(target) {}
LogicalResult
matchAndRewrite(DsAsyncBarrierArriveOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
- if (chipset < kGfx1250)
+ if (!target.has(llvm::AMDGPU::FEAT_GFX1250_INSTS))
return op->emitOpError("only supported on gfx1250+");
Location loc = op.getLoc();
@@ -3526,16 +3516,16 @@ struct DsAsyncBarrierArriveOpLowering
struct DsBarrierArriveOpLowering
: public ConvertOpToLLVMPattern<DsBarrierArriveOp> {
- Chipset chipset;
+ ROCDL::TargetInfo target;
- DsBarrierArriveOpLowering(const LLVMTypeConverter &converter, Chipset chipset)
- : ConvertOpToLLVMPattern<DsBarrierArriveOp>(converter), chipset(chipset) {
- }
+ DsBarrierArriveOpLowering(const LLVMTypeConverter &converter,
+ const ROCDL::TargetInfo &target)
+ : ConvertOpToLLVMPattern<DsBarrierArriveOp>(converter), target(target) {}
LogicalResult
matchAndRewrite(DsBarrierArriveOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
- if (chipset < kGfx1250)
+ if (!target.has(llvm::AMDGPU::FEAT_GFX1250_INSTS))
return op->emitOpError("only supported on gfx1250+");
Location loc = op.getLoc();
@@ -3668,14 +3658,15 @@ struct AMDGPUMakeDmaBaseLowering : public ConvertOpToLLVMPattern<BaseOp> {
using ConvertOpToLLVMPattern<BaseOp>::ConvertOpToLLVMPattern;
using Adaptor = typename ConvertOpToLLVMPattern<BaseOp>::OpAdaptor;
- AMDGPUMakeDmaBaseLowering(const LLVMTypeConverter &converter, Chipset chipset)
- : ConvertOpToLLVMPattern<BaseOp>(converter), chipset(chipset) {}
- Chipset chipset;
+ AMDGPUMakeDmaBaseLowering(const LLVMTypeConverter &converter,
+ const ROCDL::TargetInfo &target)
+ : ConvertOpToLLVMPattern<BaseOp>(converter), target(target) {}
+ ROCDL::TargetInfo target;
LogicalResult
matchAndRewrite(BaseOp op, Adaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
- if (chipset < kGfx1250)
+ if (!target.has(llvm::AMDGPU::FEAT_GFX1250_INSTS))
return op->emitOpError("make_dma_base is only supported on gfx1250");
Location loc = op.getLoc();
@@ -3758,9 +3749,10 @@ struct AMDGPULowerDescriptor : public ConvertOpToLLVMPattern<DescriptorOp> {
using ConvertOpToLLVMPattern<DescriptorOp>::ConvertOpToLLVMPattern;
using OpAdaptor = typename ConvertOpToLLVMPattern<DescriptorOp>::OpAdaptor;
- AMDGPULowerDescriptor(const LLVMTypeConverter &converter, Chipset chipset)
- : ConvertOpToLLVMPattern<DescriptorOp>(converter), chipset(chipset) {}
- Chipset chipset;
+ AMDGPULowerDescriptor(const LLVMTypeConverter &converter,
+ const ROCDL::TargetInfo &target)
+ : ConvertOpToLLVMPattern<DescriptorOp>(converter), target(target) {}
+ ROCDL::TargetInfo target;
Value getDGroup0(OpAdaptor &adaptor) const { return adaptor.getBase(); }
@@ -4443,7 +4435,7 @@ struct AMDGPULowerDescriptor : public ConvertOpToLLVMPattern<DescriptorOp> {
LogicalResult
matchAndRewrite(DescriptorOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
- if (chipset < kGfx1250)
+ if (!target.has(llvm::AMDGPU::FEAT_GFX1250_INSTS))
return op->emitOpError(
"make_dma_descriptor is only supported on gfx1250");
@@ -4469,14 +4461,14 @@ struct AMDGPUTensorLoadStoreOpLowering
using ConvertOpToLLVMPattern<SourceOp>::ConvertOpToLLVMPattern;
using Adaptor = typename ConvertOpToLLVMPattern<SourceOp>::OneToNOpAdaptor;
AMDGPUTensorLoadStoreOpLowering(const LLVMTypeConverter &converter,
- Chipset chipset)
- : ConvertOpToLLVMPattern<SourceOp>(converter), chipset(chipset) {}
- Chipset chipset;
+ const ROCDL::TargetInfo &target)
+ : ConvertOpToLLVMPattern<SourceOp>(converter), target(target) {}
+ ROCDL::TargetInfo target;
LogicalResult
matchAndRewrite(SourceOp op, Adaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
- if (chipset < kGfx1250)
+ if (!target.has(llvm::AMDGPU::FEAT_GFX1250_INSTS))
return op->emitOpError("is only supported on gfx1250");
ValueRange desc = adaptor.getDesc();
@@ -4496,13 +4488,14 @@ struct AMDGPUTensorLoadStoreOpLowering
struct GlobalPrefetchOpLowering
: public ConvertOpToLLVMPattern<GlobalPrefetchOp> {
- GlobalPrefetchOpLowering(const LLVMTypeConverter &converter, Chipset chipset)
- : ConvertOpToLLVMPattern<GlobalPrefetchOp>(converter), chipset(chipset) {}
+ GlobalPrefetchOpLowering(const LLVMTypeConverter &converter,
+ const ROCDL::TargetInfo &target)
+ : ConvertOpToLLVMPattern<GlobalPrefetchOp>(converter), target(target) {}
LogicalResult
matchAndRewrite(GlobalPrefetchOp op, GlobalPrefetchOpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
- if (chipset < kGfx1250)
+ if (!target.has(llvm::AMDGPU::FEAT_GFX1250_INSTS))
return op->emitOpError("is only supported on gfx1250+");
const bool isSpeculative = op.getSpeculative();
@@ -4532,7 +4525,7 @@ struct GlobalPrefetchOpLowering
}
private:
- Chipset chipset;
+ ROCDL::TargetInfo target;
};
struct ConvertAMDGPUToROCDLPass
@@ -4541,16 +4534,16 @@ struct ConvertAMDGPUToROCDLPass
void runOnOperation() override {
MLIRContext *ctx = &getContext();
- FailureOr<Chipset> maybeChipset = Chipset::parse(chipset);
- if (failed(maybeChipset)) {
- emitError(UnknownLoc::get(ctx), "Invalid chipset name: " + chipset);
+ FailureOr<ROCDL::TargetInfo> targetInfo = ROCDL::TargetInfo::get(
+ ROCDL::resolveArchOption(arch, chipset),
+ /*waveSize=*/0, [&] { return emitError(UnknownLoc::get(ctx)); });
+ if (failed(targetInfo))
return signalPassFailure();
- }
RewritePatternSet patterns(ctx);
LLVMTypeConverter converter(ctx);
- populateAMDGPUToROCDLConversionPatterns(converter, patterns, *maybeChipset);
+ populateAMDGPUToROCDLConversionPatterns(converter, patterns, *targetInfo);
amdgpu::populateCommonGPUTypeAndAttributeConversions(converter);
LLVMConversionTarget target(getContext());
target.addIllegalDialect<::mlir::amdgpu::AMDGPUDialect>();
@@ -4642,9 +4635,9 @@ void mlir::populateAMDGPUTypeAndAttributeConversions(
typeConverter.addTargetMaterialization(addUnrealizedCast);
}
-void mlir::populateAMDGPUToROCDLConversionPatterns(LLVMTypeConverter &converter,
- RewritePatternSet &patterns,
- Chipset chipset) {
+void mlir::populateAMDGPUToROCDLConversionPatterns(
+ LLVMTypeConverter &converter, RewritePatternSet &patterns,
+ const ROCDL::TargetInfo &target) {
populateAMDGPUTypeAndAttributeConversions(converter);
patterns
.add<FatRawBufferCastLowering,
@@ -4679,7 +4672,7 @@ void mlir::populateAMDGPUToROCDLConversionPatterns(LLVMTypeConverter &converter,
ROCDL::TensorStoreFromLDSOp>,
DsBarrierInitOpLowering, DsBarrierPollStateOpLowering,
DsAsyncBarrierArriveOpLowering, DsBarrierArriveOpLowering,
- GlobalPrefetchOpLowering>(converter, chipset);
+ GlobalPrefetchOpLowering>(converter, target);
patterns.add<AMDGPUSwizzleBitModeLowering, DsBarrierStatePhaseOpLowering,
DsBarrierStatePendingCountOpLowering,
DsBarrierStateInitCountOpLowering,
diff --git a/mlir/lib/Conversion/ArithToAMDGPU/ArithToAMDGPU.cpp b/mlir/lib/Conversion/ArithToAMDGPU/ArithToAMDGPU.cpp
index d0714739589f5..4a5399d7bf4bd 100644
--- a/mlir/lib/Conversion/ArithToAMDGPU/ArithToAMDGPU.cpp
+++ b/mlir/lib/Conversion/ArithToAMDGPU/ArithToAMDGPU.cpp
@@ -9,11 +9,11 @@
#include "mlir/Conversion/ArithToAMDGPU/ArithToAMDGPU.h"
#include "mlir/Dialect/AMDGPU/IR/AMDGPUDialect.h"
-#include "mlir/Dialect/AMDGPU/Utils/Chipset.h"
#include "mlir/Dialect/Arith/IR/Arith.h"
#include "mlir/Dialect/Arith/Utils/Utils.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
#include "mlir/Dialect/LLVMIR/ROCDLDialect.h"
+#include "mlir/Dialect/LLVMIR/ROCDLTargetInfo.h"
#include "mlir/Dialect/Utils/IndexingUtils.h"
#include "mlir/Dialect/Vector/IR/VectorOps.h"
#include "mlir/Dialect/Vector/Utils/VectorUtils.h"
@@ -32,9 +32,6 @@ using namespace mlir;
using namespace mlir::amdgpu;
namespace {
-// Define commonly used chipsets versions for convenience.
-constexpr Chipset kGfx942 = Chipset(9, 4, 2);
-constexpr Chipset kGfx950 = Chipset(9, 5, 0);
struct ArithToAMDGPUConversionPass final
: impl::ArithToAMDGPUConversionPassBase<ArithToAMDGPUConversionPass> {
@@ -47,10 +44,10 @@ struct ArithToAMDGPUConversionPass final
struct ExtFOnFloat8RewritePattern final : OpRewritePattern<arith::ExtFOp> {
using Base::Base;
- Chipset chipset;
- ExtFOnFloat8RewritePattern(MLIRContext *ctx, Chipset chipset,
+ ROCDL::TargetInfo target;
+ ExtFOnFloat8RewritePattern(MLIRContext *ctx, const ROCDL::TargetInfo &target,
PatternBenefit benefit)
- : OpRewritePattern::OpRewritePattern(ctx, benefit), chipset(chipset) {}
+ : OpRewritePattern::OpRewritePattern(ctx, benefit), target(target) {}
LogicalResult matchAndRewrite(arith::ExtFOp op,
PatternRewriter &rewriter) const override;
@@ -59,10 +56,11 @@ struct ExtFOnFloat8RewritePattern final : OpRewritePattern<arith::ExtFOp> {
struct TruncFToFloat8RewritePattern final : OpRewritePattern<arith::TruncFOp> {
bool saturateFP8 = false;
TruncFToFloat8RewritePattern(MLIRContext *ctx, bool saturateFP8,
- Chipset chipset, PatternBenefit benefit)
+ const ROCDL::TargetInfo &target,
+ PatternBenefit benefit)
: OpRewritePattern::OpRewritePattern(ctx, benefit),
- saturateFP8(saturateFP8), chipset(chipset) {}
- Chipset chipset;
+ saturateFP8(saturateFP8), target(target) {}
+ ROCDL::TargetInfo target;
LogicalResult matchAndRewrite(arith::TruncFOp op,
PatternRewriter &rewriter) const override;
@@ -95,10 +93,10 @@ struct ScalingTruncFRewritePattern final
} // end namespace
-static bool isSupportedF8(Type elementType, Chipset chipset) {
- if (chipset == kGfx942)
+static bool isSupportedF8(Type elementType, const ROCDL::TargetInfo &target) {
+ if (target.hasFnuzFp8())
return isa<Float8E4M3FNUZType, Float8E5M2FNUZType>(elementType);
- if (hasOcpFp8(chipset))
+ if (target.hasOcpFp8())
return isa<Float8E4M3FNType, Float8E5M2Type>(elementType);
return false;
}
@@ -126,7 +124,7 @@ ExtFOnFloat8RewritePattern::matchAndRewrite(arith::ExtFOp op,
return failure();
inType = inVecType.getElementType();
}
- if (!isSupportedF8(inType, chipset))
+ if (!isSupportedF8(inType, target))
return failure();
Location loc = op.getLoc();
@@ -278,7 +276,7 @@ TruncFToFloat8RewritePattern::matchAndRewrite(arith::TruncFOp op,
// Conversion between 8-bit floats is not supported with truncation enabled.
return failure();
- if (!isSupportedF8(outType, chipset))
+ if (!isSupportedF8(outType, target))
return failure();
Location loc = op.getLoc();
@@ -704,13 +702,13 @@ ScalingTruncFRewritePattern::matchAndRewrite(arith::ScalingTruncFOp op,
void mlir::arith::populateArithToAMDGPUConversionPatterns(
RewritePatternSet &patterns, bool convertFP8Arithmetic,
bool saturateFP8Truncf, bool allowPackedF16Rtz, bool supportsScaledExtTrunc,
- Chipset chipset, PatternBenefit benefit) {
+ const ROCDL::TargetInfo &target, PatternBenefit benefit) {
if (convertFP8Arithmetic) {
- patterns.add<ExtFOnFloat8RewritePattern>(patterns.getContext(), chipset,
+ patterns.add<ExtFOnFloat8RewritePattern>(patterns.getContext(), target,
benefit);
patterns.add<TruncFToFloat8RewritePattern>(
- patterns.getContext(), saturateFP8Truncf, chipset, benefit);
+ patterns.getContext(), saturateFP8Truncf, target, benefit);
}
if (allowPackedF16Rtz)
patterns.add<TruncfToFloat16RewritePattern>(patterns.getContext(), benefit);
@@ -725,18 +723,20 @@ void ArithToAMDGPUConversionPass::runOnOperation() {
Operation *op = getOperation();
MLIRContext *ctx = &getContext();
RewritePatternSet patterns(op->getContext());
- FailureOr<amdgpu::Chipset> maybeChipset = amdgpu::Chipset::parse(chipset);
- if (failed(maybeChipset)) {
- emitError(UnknownLoc::get(ctx), "Invalid chipset name: " + chipset);
+ FailureOr<ROCDL::TargetInfo> targetInfo = ROCDL::TargetInfo::get(
+ ROCDL::resolveArchOption(arch, chipset), /*waveSize=*/0,
+ [&] { return emitError(UnknownLoc::get(ctx)); });
+ if (failed(targetInfo)) {
return signalPassFailure();
}
bool convertFP8Arithmetic =
- *maybeChipset == kGfx942 || hasOcpFp8(*maybeChipset);
- bool supportsScaledExtTrunc = *maybeChipset == kGfx950;
+ targetInfo->has(llvm::AMDGPU::FEAT_FP8_CONVERSION_INSTS);
+ bool supportsScaledExtTrunc =
+ targetInfo->has(llvm::AMDGPU::FEAT_GFX950_INSTS);
arith::populateArithToAMDGPUConversionPatterns(
patterns, convertFP8Arithmetic, saturateFP8Truncf, allowPackedF16Rtz,
- supportsScaledExtTrunc, *maybeChipset);
+ supportsScaledExtTrunc, *targetInfo);
if (failed(applyPatternsGreedily(op, std::move(patterns))))
return signalPassFailure();
}
diff --git a/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp b/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp
index c6f81fd533dff..94f0f82b2c39e 100644
--- a/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp
+++ b/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp
@@ -241,18 +241,18 @@ struct GPUSubgroupSizeOpToROCDL : ConvertOpToLLVMPattern<gpu::SubgroupSizeOp> {
using ConvertOpToLLVMPattern::ConvertOpToLLVMPattern;
GPUSubgroupSizeOpToROCDL(const LLVMTypeConverter &converter,
- amdgpu::Chipset chipset)
- : ConvertOpToLLVMPattern<gpu::SubgroupSizeOp>(converter),
- chipset(chipset) {}
+ const ROCDL::TargetInfo &target)
+ : ConvertOpToLLVMPattern<gpu::SubgroupSizeOp>(converter), target(target) {
+ }
LogicalResult
matchAndRewrite(gpu::SubgroupSizeOp op, gpu::SubgroupSizeOp::Adaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
LLVM::ConstantRangeAttr bounds = nullptr;
- bool isBeforeGfx10 = chipset.majorVersion < 10;
+ bool isWave64 = target.getWavefrontSize() == 64;
if (auto upperBoundAttr = op.getUpperBoundAttr()) {
bounds = rewriter.getAttr<LLVM::ConstantRangeAttr>(
- /*bitWidth=*/32, /*lower=*/isBeforeGfx10 ? 64 : 32,
+ /*bitWidth=*/32, /*lower=*/isWave64 ? 64 : 32,
/*upper=*/op.getUpperBoundAttr().getInt() + 1);
}
Value wavefrontOp = ROCDL::WavefrontSizeOp::create(
@@ -263,16 +263,15 @@ struct GPUSubgroupSizeOpToROCDL : ConvertOpToLLVMPattern<gpu::SubgroupSizeOp> {
return success();
}
- const amdgpu::Chipset chipset;
+ const ROCDL::TargetInfo target;
};
struct GPUSubgroupIdOpToROCDL : ConvertOpToLLVMPattern<gpu::SubgroupIdOp> {
using ConvertOpToLLVMPattern::ConvertOpToLLVMPattern;
GPUSubgroupIdOpToROCDL(const LLVMTypeConverter &converter,
- amdgpu::Chipset chipset)
- : ConvertOpToLLVMPattern<gpu::SubgroupIdOp>(converter), chipset(chipset) {
- }
+ const ROCDL::TargetInfo &target)
+ : ConvertOpToLLVMPattern<gpu::SubgroupIdOp>(converter), target(target) {}
LogicalResult
matchAndRewrite(gpu::SubgroupIdOp op, gpu::SubgroupIdOp::Adaptor adaptor,
@@ -281,7 +280,7 @@ struct GPUSubgroupIdOpToROCDL : ConvertOpToLLVMPattern<gpu::SubgroupIdOp> {
auto int32Type = rewriter.getI32Type();
Value subgroupId;
- if (chipset.majorVersion >= 12) {
+ if (target.has(llvm::AMDGPU::FEAT_GFX12_INSTS)) {
// For gfx12+, use the hardware wave.id register directly.
LLVM::ConstantRangeAttr bounds;
if (auto upperBoundAttr = op.getUpperBoundAttr())
@@ -345,7 +344,7 @@ struct GPUSubgroupIdOpToROCDL : ConvertOpToLLVMPattern<gpu::SubgroupIdOp> {
return success();
}
- const amdgpu::Chipset chipset;
+ const ROCDL::TargetInfo target;
};
static bool isSupportedReadLaneType(Type type) {
@@ -567,10 +566,10 @@ static constexpr int32_t kWholeClusterBarrierId = -3;
static constexpr int32_t kWholeWorkgroupBarrierId = -1;
struct GPUBarrierOpLowering final : ConvertOpToLLVMPattern<gpu::BarrierOp> {
GPUBarrierOpLowering(const LLVMTypeConverter &converter,
- amdgpu::Chipset chipset)
- : ConvertOpToLLVMPattern<gpu::BarrierOp>(converter), chipset(chipset) {}
+ const ROCDL::TargetInfo &target)
+ : ConvertOpToLLVMPattern<gpu::BarrierOp>(converter), target(target) {}
- amdgpu::Chipset chipset;
+ ROCDL::TargetInfo target;
LogicalResult
matchAndRewrite(gpu::BarrierOp op, gpu::BarrierOp::Adaptor adaptor,
@@ -591,7 +590,7 @@ struct GPUBarrierOpLowering final : ConvertOpToLLVMPattern<gpu::BarrierOp> {
// Cluster scope: gfx1250+ only, signal/wait with constant -3.
if (scope == gpu::BarrierScope::Cluster) {
- if (chipset < amdgpu::Chipset(12, 5, 0))
+ if (!target.has(llvm::AMDGPU::FEAT_GFX1250_INSTS))
return op.emitOpError("cluster scope barriers require gfx1250+");
emitFences(op.getAddressSpaces(), rewriter, loc, "cluster",
/*before=*/true);
@@ -609,7 +608,7 @@ struct GPUBarrierOpLowering final : ConvertOpToLLVMPattern<gpu::BarrierOp> {
// Named barrier path.
if (Value namedBarrier = adaptor.getNamedBarrier()) {
- if (chipset.majorVersion < 12)
+ if (!target.has(llvm::AMDGPU::FEAT_GFX12_INSTS))
return op.emitOpError("named barriers require gfx12+");
emitFences(op.getAddressSpaces(), rewriter, loc, "workgroup",
@@ -631,7 +630,7 @@ struct GPUBarrierOpLowering final : ConvertOpToLLVMPattern<gpu::BarrierOp> {
// Regular workgroup barrier.
emitFences(op.getAddressSpaces(), rewriter, loc, "workgroup",
/*before=*/true);
- if (chipset.majorVersion < 12) {
+ if (!target.has(llvm::AMDGPU::FEAT_GFX12_INSTS)) {
ROCDL::SBarrierOp::create(rewriter, loc);
} else {
ROCDL::BarrierSignalOp::create(rewriter, loc, kWholeWorkgroupBarrierId);
@@ -648,17 +647,17 @@ struct GPUBarrierOpLowering final : ConvertOpToLLVMPattern<gpu::BarrierOp> {
struct GPUInitializeNamedBarrierOpLowering final
: ConvertOpToLLVMPattern<gpu::InitializeNamedBarrierOp> {
GPUInitializeNamedBarrierOpLowering(const LLVMTypeConverter &converter,
- amdgpu::Chipset chipset)
+ const ROCDL::TargetInfo &target)
: ConvertOpToLLVMPattern<gpu::InitializeNamedBarrierOp>(converter),
- chipset(chipset) {}
+ target(target) {}
- amdgpu::Chipset chipset;
+ ROCDL::TargetInfo target;
LogicalResult
matchAndRewrite(gpu::InitializeNamedBarrierOp op,
gpu::InitializeNamedBarrierOp::Adaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
- if (chipset.majorVersion < 12)
+ if (!target.has(llvm::AMDGPU::FEAT_GFX12_INSTS))
return op.emitOpError("named barriers require gfx12+");
Location loc = op.getLoc();
@@ -757,9 +756,10 @@ struct LowerGpuOpsToROCDLOpsPass final
UnitAttr::get(ctx));
}
- FailureOr<amdgpu::Chipset> maybeChipset = amdgpu::Chipset::parse(chipset);
- if (failed(maybeChipset)) {
- emitError(UnknownLoc::get(ctx), "Invalid chipset name: " + chipset);
+ FailureOr<ROCDL::TargetInfo> targetInfo = ROCDL::TargetInfo::get(
+ ROCDL::resolveArchOption(arch, chipset), waveSize,
+ [&] { return emitError(UnknownLoc::get(ctx)); });
+ if (failed(targetInfo)) {
return signalPassFailure();
}
@@ -792,7 +792,7 @@ struct LowerGpuOpsToROCDLOpsPass final
{
RewritePatternSet patterns(ctx);
populateGpuRewritePatterns(patterns);
- populateGpuPromoteShuffleToAMDGPUPatterns(patterns, maybeChipset);
+ populateGpuPromoteShuffleToAMDGPUPatterns(patterns, *targetInfo);
(void)applyPatternsGreedily(m, std::move(patterns));
}
@@ -828,9 +828,9 @@ struct LowerGpuOpsToROCDLOpsPass final
}
populateAMDGPUToROCDLConversionPatterns(converter, llvmPatterns,
- *maybeChipset);
+ *targetInfo);
populateGpuToROCDLConversionPatterns(converter, llvmPatterns, runtime,
- *maybeChipset);
+ *targetInfo);
configureGpuToROCDLConversionLegality(target);
if (failed(applyPartialConversion(m, target, std::move(llvmPatterns))))
signalPassFailure();
@@ -878,7 +878,7 @@ void mlir::configureGpuToROCDLConversionLegality(ConversionTarget &target) {
void mlir::populateGpuToROCDLConversionPatterns(
const LLVMTypeConverter &converter, RewritePatternSet &patterns,
- mlir::gpu::amd::Runtime runtime, amdgpu::Chipset chipset) {
+ mlir::gpu::amd::Runtime runtime, const ROCDL::TargetInfo &target) {
using gpu::index_lowering::IndexKind;
using gpu::index_lowering::IntrType;
using mlir::gpu::amd::Runtime;
@@ -917,7 +917,7 @@ void mlir::populateGpuToROCDLConversionPatterns(
GPUSubgroupBroadcastOpToROCDL, GPUBallotOpToROCDL>(converter);
patterns.add<GPUSubgroupIdOpToROCDL, GPUSubgroupSizeOpToROCDL,
GPUBarrierOpLowering, GPUInitializeNamedBarrierOpLowering>(
- converter, chipset);
+ converter, target);
- populateMathToROCDLConversionPatterns(converter, patterns, chipset);
+ populateMathToROCDLConversionPatterns(converter, patterns, target);
}
diff --git a/mlir/lib/Conversion/MathToROCDL/MathToROCDL.cpp b/mlir/lib/Conversion/MathToROCDL/MathToROCDL.cpp
index 04c3c9e4136e7..b54de7e87adab 100644
--- a/mlir/lib/Conversion/MathToROCDL/MathToROCDL.cpp
+++ b/mlir/lib/Conversion/MathToROCDL/MathToROCDL.cpp
@@ -11,11 +11,11 @@
#include "mlir/Conversion/LLVMCommon/LoweringOptions.h"
#include "mlir/Conversion/LLVMCommon/TypeConverter.h"
#include "mlir/Conversion/LLVMCommon/VectorPattern.h"
-#include "mlir/Dialect/AMDGPU/Utils/Chipset.h"
#include "mlir/Dialect/Arith/IR/Arith.h"
#include "mlir/Dialect/Func/IR/FuncDialect.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
#include "mlir/Dialect/LLVMIR/ROCDLDialect.h"
+#include "mlir/Dialect/LLVMIR/ROCDLTargetInfo.h"
#include "mlir/Dialect/Math/IR/Math.h"
#include "mlir/Dialect/Vector/IR/VectorOps.h"
#include "mlir/IR/BuiltinDialect.h"
@@ -85,7 +85,7 @@ struct ClampFOpConversion final
void mlir::populateMathToROCDLConversionPatterns(
const LLVMTypeConverter &converter, RewritePatternSet &patterns,
- std::optional<amdgpu::Chipset> chipset) {
+ std::optional<ROCDL::TargetInfo> target) {
// Handled by mathToLLVM: math::AbsIOp
// Handled by mathToLLVM: math::AbsFOp
// Handled by mathToLLVM: math::CopySignOp
@@ -161,10 +161,10 @@ void mlir::populateMathToROCDLConversionPatterns(
populateOpPatterns<arith::RemFOp>(converter, patterns, "__ocml_fmod_f32",
"__ocml_fmod_f64", "__ocml_fmod_f16");
- if (chipset.has_value() && chipset->majorVersion >= 9) {
+ if (target && target->has(llvm::AMDGPU::FEAT_GFX9_INSTS)) {
patterns.add<ClampFOpConversion>(converter);
} else {
- LDBG() << "Chipset dependent patterns were not added";
+ LDBG() << "Target dependent patterns were not added";
}
}
@@ -184,15 +184,20 @@ void ConvertMathToROCDLPass::runOnOperation() {
LowerToLLVMOptions options(ctx, DataLayout(m));
LLVMTypeConverter converter(ctx, options);
- FailureOr<amdgpu::Chipset> maybeChipset;
- if (!chipset.empty()) {
- maybeChipset = amdgpu::Chipset::parse(chipset);
- if (failed(maybeChipset))
+ // An empty architecture means "no target", in which case the
+ // target-dependent patterns are simply not added.
+ std::optional<ROCDL::TargetInfo> resolved;
+ StringRef resolvedArch = ROCDL::resolveArchOption(arch, chipset);
+ if (!resolvedArch.empty()) {
+ FailureOr<ROCDL::TargetInfo> targetInfo =
+ ROCDL::TargetInfo::get(resolvedArch, /*waveSize=*/0, [&] {
+ return emitError(UnknownLoc::get(&getContext()));
+ });
+ if (failed(targetInfo))
return signalPassFailure();
+ resolved = *targetInfo;
}
- populateMathToROCDLConversionPatterns(
- converter, patterns,
- succeeded(maybeChipset) ? std::optional(*maybeChipset) : std::nullopt);
+ populateMathToROCDLConversionPatterns(converter, patterns, resolved);
ConversionTarget target(getContext());
target
diff --git a/mlir/lib/Dialect/AMDGPU/Transforms/CMakeLists.txt b/mlir/lib/Dialect/AMDGPU/Transforms/CMakeLists.txt
index 29baef635ec80..28fa958181d03 100644
--- a/mlir/lib/Dialect/AMDGPU/Transforms/CMakeLists.txt
+++ b/mlir/lib/Dialect/AMDGPU/Transforms/CMakeLists.txt
@@ -16,6 +16,7 @@ add_mlir_dialect_library(MLIRAMDGPUTransforms
MLIRAffineUtils
MLIRArithDialect
MLIRMemRefDialect
+ MLIRROCDLDialect
MLIRSCFDialect
MLIRVectorDialect
MLIRControlFlowDialect
diff --git a/mlir/lib/Dialect/AMDGPU/Transforms/EmulateAtomics.cpp b/mlir/lib/Dialect/AMDGPU/Transforms/EmulateAtomics.cpp
index ce47da8d9ee97..930235995b4e1 100644
--- a/mlir/lib/Dialect/AMDGPU/Transforms/EmulateAtomics.cpp
+++ b/mlir/lib/Dialect/AMDGPU/Transforms/EmulateAtomics.cpp
@@ -9,9 +9,9 @@
#include "mlir/Dialect/AMDGPU/Transforms/Passes.h"
#include "mlir/Dialect/AMDGPU/IR/AMDGPUDialect.h"
-#include "mlir/Dialect/AMDGPU/Utils/Chipset.h"
#include "mlir/Dialect/Arith/IR/Arith.h"
#include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h"
+#include "mlir/Dialect/LLVMIR/ROCDLTargetInfo.h"
#include "mlir/Dialect/Vector/IR/VectorOps.h"
#include "mlir/IR/BuiltinAttributes.h"
#include "mlir/IR/TypeUtilities.h"
@@ -193,43 +193,49 @@ LogicalResult RawBufferAtomicByCasPattern<AtomicOp, ArithOp>::matchAndRewrite(
return success();
}
+/// Returns whether \p op can be lowered to a native buffer atomic fadd on
+/// \p target.
+static bool isFaddNativelySupported(const ROCDL::TargetInfo &target,
+ RawBufferAtomicFaddOp op) {
+ namespace AMDGPU = ::llvm::AMDGPU;
+
+ // A target with only the no-return form can still perform the atomic, it
+ // just cannot report the old value, so it suffices when the result is dead.
+ bool hasFadd = target.has(AMDGPU::FEAT_ATOMIC_FADD_RTN_INSTS) ||
+ (target.has(AMDGPU::FEAT_ATOMIC_FADD_NO_RTN_INSTS) &&
+ op.getOldValue().use_empty());
+ if (!hasFadd)
+ return false;
+
+ // The packed 16-bit forms are separate instructions with their own features.
+ Type elemType = getElementTypeOrSelf(op.getValue().getType());
+ if (isa<Float16Type>(elemType))
+ return target.has(AMDGPU::FEAT_ATOMIC_BUFFER_GLOBAL_PK_ADD_F16_INSTS);
+ if (isa<BFloat16Type>(elemType))
+ return target.has(AMDGPU::FEAT_ATOMIC_BUFFER_PK_ADD_BF16_INST);
+ return true;
+}
+
void mlir::amdgpu::populateAmdgpuEmulateAtomicsPatterns(
- ConversionTarget &target, RewritePatternSet &patterns, Chipset chipset,
- PatternBenefit benefit) {
- // gfx10 has no atomic adds.
- if (chipset.majorVersion == 10 || chipset < Chipset(9, 0, 8)) {
- target.addIllegalOp<RawBufferAtomicFaddOp>();
- }
- // gfx11 has no fp16 atomics
- if (chipset.majorVersion == 11) {
- target.addDynamicallyLegalOp<RawBufferAtomicFaddOp>(
- [](RawBufferAtomicFaddOp op) -> bool {
- Type elemType = getElementTypeOrSelf(op.getValue().getType());
- return !isa<Float16Type, BFloat16Type>(elemType);
+ ConversionTarget &target, RewritePatternSet &patterns,
+ const ROCDL::TargetInfo &targetInfo, PatternBenefit benefit) {
+ namespace AMDGPU = ::llvm::AMDGPU;
+
+ target.addDynamicallyLegalOp<RawBufferAtomicFaddOp>(
+ [targetInfo](RawBufferAtomicFaddOp op) -> bool {
+ return isFaddNativelySupported(targetInfo, op);
+ });
+
+ // Floating-point min and max are only emulated on gfx9; later generations
+ // have them for every type this op accepts. f64 is the one gfx9 case that
+ // may be native.
+ if (targetInfo.isGeneration(9)) {
+ target.addDynamicallyLegalOp<RawBufferAtomicFmaxOp>(
+ [targetInfo](RawBufferAtomicFmaxOp op) -> bool {
+ return op.getValue().getType().isF64() &&
+ targetInfo.has(AMDGPU::FEAT_ATOMIC_FMIN_FMAX_GLOBAL_F64);
});
}
- // gfx9 has no to a very limited support for floating-point min and max.
- if (chipset.majorVersion == 9) {
- if (chipset >= Chipset(9, 0, 0xa)) {
- // gfx90a supports f64 max (and min, but we don't have a min wrapper right
- // now) but all other types need to be emulated.
- target.addDynamicallyLegalOp<RawBufferAtomicFmaxOp>(
- [](RawBufferAtomicFmaxOp op) -> bool {
- return op.getValue().getType().isF64();
- });
- } else {
- target.addIllegalOp<RawBufferAtomicFmaxOp>();
- }
- // TODO(https://github.com/llvm/llvm-project/issues/129206): Refactor
- // this to avoid hardcoding ISA version: gfx950 has bf16 atomics.
- if (chipset < Chipset(9, 5, 0)) {
- target.addDynamicallyLegalOp<RawBufferAtomicFaddOp>(
- [](RawBufferAtomicFaddOp op) -> bool {
- Type elemType = getElementTypeOrSelf(op.getValue().getType());
- return !isa<BFloat16Type>(elemType);
- });
- }
- }
patterns.add<
RawBufferAtomicByCasPattern<RawBufferAtomicFaddOp, arith::AddFOp>,
RawBufferAtomicByCasPattern<RawBufferAtomicFmaxOp, arith::MaximumFOp>,
@@ -240,11 +246,11 @@ void mlir::amdgpu::populateAmdgpuEmulateAtomicsPatterns(
void AmdgpuEmulateAtomicsPass::runOnOperation() {
Operation *op = getOperation();
- FailureOr<Chipset> maybeChipset = Chipset::parse(chipset);
- if (failed(maybeChipset)) {
- emitError(op->getLoc(), "Invalid chipset name: " + chipset);
+ FailureOr<ROCDL::TargetInfo> targetInfo =
+ ROCDL::TargetInfo::get(ROCDL::resolveArchOption(arch, chipset),
+ /*waveSize=*/0, [&] { return op->emitError(); });
+ if (failed(targetInfo))
return signalPassFailure();
- }
MLIRContext &ctx = getContext();
ConversionTarget target(ctx);
@@ -252,7 +258,7 @@ void AmdgpuEmulateAtomicsPass::runOnOperation() {
target.markUnknownOpDynamicallyLegal(
[](Operation *op) -> bool { return true; });
- populateAmdgpuEmulateAtomicsPatterns(target, patterns, *maybeChipset);
+ populateAmdgpuEmulateAtomicsPatterns(target, patterns, *targetInfo);
if (failed(applyPartialConversion(op, target, std::move(patterns))))
return signalPassFailure();
}
diff --git a/mlir/lib/Dialect/GPU/Pipelines/GPUToROCDLPipeline.cpp b/mlir/lib/Dialect/GPU/Pipelines/GPUToROCDLPipeline.cpp
index 1e5fd09a00a75..edbcaa1f8ca1c 100644
--- a/mlir/lib/Dialect/GPU/Pipelines/GPUToROCDLPipeline.cpp
+++ b/mlir/lib/Dialect/GPU/Pipelines/GPUToROCDLPipeline.cpp
@@ -28,6 +28,7 @@
#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/GPU/Pipelines/Passes.h"
#include "mlir/Dialect/GPU/Transforms/Passes.h"
+#include "mlir/Dialect/LLVMIR/ROCDLTargetInfo.h"
#include "mlir/Dialect/MemRef/Transforms/Passes.h"
#include "mlir/Pass/PassManager.h"
#include "mlir/Pass/PassOptions.h"
@@ -42,12 +43,13 @@ namespace {
//===----------------------------------------------------------------------===//
void buildCommonPassPipeline(
OpPassManager &pm, const mlir::gpu::GPUToROCDLPipelineOptions &options) {
+ std::string arch = ROCDL::resolveArchOption(options.arch, options.chip).str();
// Lower AMDGPU dialect ops (e.g. amdgpu.lds_barrier, amdgpu.dpp,
// amdgpu.mfma, amdgpu.dot, ...) to ROCDL intrinsics first, while they may
// still live in unout-lined `gpu.launch` bodies. Mirrors the way NVVM's
// pipeline runs `convert-nvgpu-to-nvvm` before kernel outlining.
ConvertAMDGPUToROCDLPassOptions amdgpuToROCDLOpt;
- amdgpuToROCDLOpt.chipset = options.chip;
+ amdgpuToROCDLOpt.arch = arch;
pm.addPass(createConvertAMDGPUToROCDLPass(amdgpuToROCDLOpt));
pm.addPass(createGpuKernelOutliningPass());
@@ -57,12 +59,10 @@ void buildCommonPassPipeline(
pm.addPass(memref::createExpandStridedMetadataPass());
GpuROCDLAttachTargetOptions rocdlTargetOptions;
- rocdlTargetOptions.triple = options.triple;
- rocdlTargetOptions.chip = options.chip;
- rocdlTargetOptions.features = options.features;
+ rocdlTargetOptions.arch = arch;
rocdlTargetOptions.abiVersion = options.abiVersion;
rocdlTargetOptions.optLevel = options.optLevel;
- rocdlTargetOptions.wave64Flag = options.wave64;
+ rocdlTargetOptions.waveSize = options.waveSize;
pm.addPass(createGpuROCDLAttachTarget(rocdlTargetOptions));
pm.addPass(createLowerAffinePass());
@@ -79,8 +79,10 @@ void buildCommonPassPipeline(
//===----------------------------------------------------------------------===//
void buildGpuPassPipeline(OpPassManager &pm,
const mlir::gpu::GPUToROCDLPipelineOptions &options) {
+ std::string arch = ROCDL::resolveArchOption(options.arch, options.chip).str();
ConvertGpuOpsToROCDLOpsOptions opt;
- opt.chipset = options.chip;
+ opt.arch = arch;
+ opt.waveSize = options.waveSize;
opt.useBarePtrCallConv = options.kernelUseBarePtrCallConv;
opt.indexBitwidth = options.indexBitWidth;
// Always declare HIP as the runtime so that gpu.printf etc. lower to the
diff --git a/mlir/lib/Dialect/GPU/TransformOps/GPUTransformOps.cpp b/mlir/lib/Dialect/GPU/TransformOps/GPUTransformOps.cpp
index 9591b76a5330c..a7858781a1079 100644
--- a/mlir/lib/Dialect/GPU/TransformOps/GPUTransformOps.cpp
+++ b/mlir/lib/Dialect/GPU/TransformOps/GPUTransformOps.cpp
@@ -14,7 +14,6 @@
#include "mlir/Conversion/GPUToROCDL/GPUToROCDLPass.h"
#include "mlir/Conversion/LLVMCommon/TypeConverter.h"
#include "mlir/Dialect/AMDGPU/IR/AMDGPUDialect.h"
-#include "mlir/Dialect/AMDGPU/Utils/Chipset.h"
#include "mlir/Dialect/Arith/IR/Arith.h"
#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/GPU/TransformOps/Utils.h"
@@ -108,21 +107,20 @@ void transform::ApplyGPUToROCDLConversionPatternsOp::populatePatterns(
TypeConverter &typeConverter, RewritePatternSet &patterns) {
auto &llvmTypeConverter = static_cast<LLVMTypeConverter &>(typeConverter);
amdgpu::populateCommonGPUTypeAndAttributeConversions(llvmTypeConverter);
- FailureOr<amdgpu::Chipset> maybeChipset =
- amdgpu::Chipset::parse(getChipset());
- assert(llvm::succeeded(maybeChipset) && "expected valid chipset");
+ // The verifier has already rejected anything unparseable.
+ FailureOr<ROCDL::TargetInfo> targetInfo =
+ ROCDL::TargetInfo::get(getArch(), getWavesize().value_or(0));
+ assert(llvm::succeeded(targetInfo) && "verifier accepted this target");
populateGpuToROCDLConversionPatterns(
- llvmTypeConverter, patterns, mlir::gpu::amd::Runtime::HIP, *maybeChipset);
+ llvmTypeConverter, patterns, mlir::gpu::amd::Runtime::HIP, *targetInfo);
}
LogicalResult
transform::ApplyGPUToROCDLConversionPatternsOp::verifyTypeConverter(
transform::TypeConverterBuilderOpInterface builder) {
- FailureOr<amdgpu::Chipset> maybeChipset =
- amdgpu::Chipset::parse(getChipset());
- if (failed(maybeChipset)) {
- return emitOpError("Invalid chipset name: " + getChipset());
- }
+ if (failed(ROCDL::TargetInfo::get(getArch(), getWavesize().value_or(0),
+ [&] { return emitOpError(); })))
+ return failure();
if (builder.getTypeConverterType() != "LLVMTypeConverter")
return emitOpError("expected LLVMTypeConverter");
return success();
@@ -138,16 +136,27 @@ void ApplyGPURewritePatternsOp::populatePatterns(RewritePatternSet &patterns) {
void transform::ApplyGPUPromoteShuffleToAMDGPUPatternsOp::populatePatterns(
RewritePatternSet &patterns) {
- std::optional<StringRef> chipsetName = getChipset();
- std::optional<amdgpu::Chipset> maybeChipset;
- if (chipsetName) {
- FailureOr<amdgpu::Chipset> parsedChipset =
- amdgpu::Chipset::parse(*chipsetName);
- assert(llvm::succeeded(parsedChipset) && "expected valid chipset");
- maybeChipset = parsedChipset;
+ std::optional<StringRef> archName = getArch();
+ std::optional<ROCDL::TargetInfo> targetInfo;
+ if (archName) {
+ // The verifier has already rejected anything unparseable.
+ FailureOr<ROCDL::TargetInfo> parsed = ROCDL::TargetInfo::get(*archName);
+ assert(llvm::succeeded(parsed) && "verifier accepted this target");
+ targetInfo = *parsed;
}
- populateGpuPromoteShuffleToAMDGPUPatterns(patterns, maybeChipset);
+ populateGpuPromoteShuffleToAMDGPUPatterns(patterns, targetInfo);
+}
+
+LogicalResult transform::ApplyGPUPromoteShuffleToAMDGPUPatternsOp::verify() {
+ std::optional<StringRef> archName = getArch();
+ if (!archName)
+ return success();
+
+ if (failed(ROCDL::TargetInfo::get(*archName, /*waveSize=*/0,
+ [&] { return emitOpError(); })))
+ return failure();
+ return success();
}
//===----------------------------------------------------------------------===//
diff --git a/mlir/lib/Dialect/GPU/Transforms/PromoteShuffleToAMDGPU.cpp b/mlir/lib/Dialect/GPU/Transforms/PromoteShuffleToAMDGPU.cpp
index 01da26f88a84d..8fb62258c98ea 100644
--- a/mlir/lib/Dialect/GPU/Transforms/PromoteShuffleToAMDGPU.cpp
+++ b/mlir/lib/Dialect/GPU/Transforms/PromoteShuffleToAMDGPU.cpp
@@ -11,8 +11,8 @@
//
//===----------------------------------------------------------------------===//
-#include "mlir/Dialect/AMDGPU/Utils/Chipset.h"
#include "mlir/Dialect/GPU/Transforms/Passes.h"
+#include "mlir/Dialect/LLVMIR/ROCDLTargetInfo.h"
#include "mlir/Dialect/AMDGPU/IR/AMDGPUDialect.h"
#include "mlir/Dialect/Arith/IR/Arith.h"
@@ -24,8 +24,6 @@ using namespace mlir;
namespace {
-constexpr amdgpu::Chipset kGfx950 = amdgpu::Chipset(9, 5, 0);
-
/// Try to promote `gpu.shuffle` to `amdgpu.swizzle_bitmode`, width must be 64
/// and offset must be a constant integer in the range [0, 31].
struct PromoteShuffleToSwizzlePattern
@@ -100,10 +98,10 @@ struct PromoteShuffleToPermlanePattern
} // namespace
void mlir::populateGpuPromoteShuffleToAMDGPUPatterns(
- RewritePatternSet &patterns, std::optional<amdgpu::Chipset> maybeChipset) {
+ RewritePatternSet &patterns, std::optional<ROCDL::TargetInfo> target) {
patterns.add<PromoteShuffleToSwizzlePattern>(patterns.getContext(),
/*benefit*/ 1);
- if (maybeChipset && *maybeChipset >= kGfx950)
+ if (target && target->has(llvm::AMDGPU::FEAT_PERMLANE32_SWAP))
patterns.add<PromoteShuffleToPermlanePattern>(patterns.getContext(),
/*benefit*/ 2);
}
diff --git a/mlir/lib/Dialect/GPU/Transforms/ROCDLAttachTarget.cpp b/mlir/lib/Dialect/GPU/Transforms/ROCDLAttachTarget.cpp
index 1f44ffa52e068..cef6deee5c354 100644
--- a/mlir/lib/Dialect/GPU/Transforms/ROCDLAttachTarget.cpp
+++ b/mlir/lib/Dialect/GPU/Transforms/ROCDLAttachTarget.cpp
@@ -15,6 +15,7 @@
#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/LLVMIR/ROCDLDialect.h"
+#include "mlir/Dialect/LLVMIR/ROCDLTargetInfo.h"
#include "mlir/IR/Builders.h"
#include "mlir/Pass/Pass.h"
#include "mlir/Target/LLVM/ROCDL/Target.h"
@@ -33,7 +34,7 @@ struct ROCDLAttachTarget
: public impl::GpuROCDLAttachTargetBase<ROCDLAttachTarget> {
using Base::Base;
- DictionaryAttr getFlags(OpBuilder &builder) const;
+ DictionaryAttr getFlags(OpBuilder &builder, bool isWave64) const;
void runOnOperation() override;
@@ -43,13 +44,14 @@ struct ROCDLAttachTarget
};
} // namespace
-DictionaryAttr ROCDLAttachTarget::getFlags(OpBuilder &builder) const {
+DictionaryAttr ROCDLAttachTarget::getFlags(OpBuilder &builder,
+ bool isWave64) const {
UnitAttr unitAttr = builder.getUnitAttr();
SmallVector<NamedAttribute, 6> flags;
auto addFlag = [&](StringRef flag) {
flags.push_back(builder.getNamedAttr(flag, unitAttr));
};
- if (!wave64Flag)
+ if (!isWave64)
addFlag("no_wave64");
if (fastFlag)
addFlag("fast");
@@ -68,10 +70,89 @@ DictionaryAttr ROCDLAttachTarget::getFlags(OpBuilder &builder) const {
void ROCDLAttachTarget::runOnOperation() {
OpBuilder builder(&getContext());
+
+ // #rocdl.target feeds the TargetMachine, whose -mcpu is a bare processor
+ // name. Split an `arch` into the pieces the attribute wants, mirroring
+ // Clang.
+ std::string resolvedTriple = triple;
+ std::string resolvedChip = chip;
+ std::string resolvedFeatures = features;
+ // `wavesize` is the only wavefront-size control. Without an `arch` there is
+ // no target to ask, so leaving it at 0 keeps the Wave64 that `#rocdl.target`
+ // has always assumed.
+ bool resolvedWave64 = waveSize != 32;
+ // Set when `arch` was given and used so its xnack/sramecc modifiers can
+ // become module flags.
+ std::optional<ROCDL::TargetInfo> targetInfo;
+ if (!arch.empty()) {
+ std::optional<llvm::AMDGPU::TargetID> id =
+ ROCDL::TargetInfo::parseTargetID(arch);
+ if (!id) {
+ emitError(UnknownLoc::get(&getContext()))
+ << "'" << arch << "' is not a valid AMDGPU architecture";
+ return signalPassFailure();
+ }
+ // #rocdl.target requires a chip, so a triple that names no GPU (the legacy
+ // subarch-less "amdgcn-amd-amdhsa") cannot be attached.
+ if (id->getGPUKind() == llvm::AMDGPU::GK_NONE) {
+ emitError(UnknownLoc::get(&getContext()))
+ << "'" << arch
+ << "' names no GPU; a chip is required to attach a "
+ "target";
+ return signalPassFailure();
+ }
+
+ llvm::Triple parsed(id->getTargetTripleString());
+
+ StringRef archName = parsed.getArchName();
+ llvm::Triple::SubArchType subArch =
+ llvm::AMDGPU::getSubArch(id->getGPUKind());
+ if (StringRef subArchName = llvm::AMDGPU::getSubArchName(subArch);
+ !subArchName.empty())
+ archName = subArchName;
+
+ // Drop the "unknown" environment part of triples since a lot of the
+ // toolchain expects a 3-component form.
+ resolvedTriple =
+ (parsed.getEnvironment() == llvm::Triple::UnknownEnvironment
+ ? llvm::Triple(archName, parsed.getVendorName(),
+ parsed.getOSName())
+ : llvm::Triple(archName, parsed.getVendorName(),
+ parsed.getOSName(), parsed.getEnvironmentName()))
+ .str();
+ resolvedChip = llvm::AMDGPU::getArchNameAMDGCN(id->getGPUKind()).str();
+
+ // Take the wavefront size from the target, since `wavesize`'s Wave64
+ // fallback is wrong on targets like gfx10.
+ FailureOr<ROCDL::TargetInfo> info =
+ ROCDL::TargetInfo::get(arch, waveSize, [&] {
+ return emitError(UnknownLoc::get(&getContext()));
+ });
+ if (failed(info))
+ return signalPassFailure();
+ targetInfo = *info;
+ resolvedWave64 = info->getWavefrontSize() == 64;
+
+ // Record the wavefrontsize option into the features set so that the device
+ // libraries and codegen agree with each other in cases where both
+ // wavefontsize32 and wavefrontsize64 are permitted options.
+ if (info->supportsBothWavefrontSizes()) {
+ if (!resolvedFeatures.empty())
+ resolvedFeatures += ",";
+ resolvedFeatures +=
+ resolvedWave64 ? "+wavefrontsize64" : "+wavefrontsize32";
+ }
+ } else if (waveSize != 0 && waveSize != 32 && waveSize != 64) {
+ emitError(UnknownLoc::get(&getContext()))
+ << "wavefront size must be 32 or 64, got " << waveSize;
+ return signalPassFailure();
+ }
+
ArrayRef<std::string> libs(linkLibs);
SmallVector<StringRef> filesToLink(libs);
auto target = builder.getAttr<ROCDLTargetAttr>(
- optLevel, triple, chip, features, abiVersion, getFlags(builder),
+ optLevel, resolvedTriple, resolvedChip, resolvedFeatures, abiVersion,
+ getFlags(builder, resolvedWave64),
filesToLink.empty() ? nullptr : builder.getStrArrayAttr(filesToLink));
llvm::Regex matcher(moduleMatcher);
for (Region ®ion : getOperation()->getRegions())
@@ -89,5 +170,7 @@ void ROCDLAttachTarget::runOnOperation() {
targets.erase(llvm::unique(targets), targets.end());
// Update the target attribute array.
module.setTargetsAttr(builder.getArrayAttr(targets));
+ if (targetInfo)
+ targetInfo->migrateArchFeaturesToModuleFlags(module);
}
}
diff --git a/mlir/lib/Dialect/GPU/Transforms/SubgroupReduceLowering.cpp b/mlir/lib/Dialect/GPU/Transforms/SubgroupReduceLowering.cpp
index ec1571a56fe4a..23f1a16ca85aa 100644
--- a/mlir/lib/Dialect/GPU/Transforms/SubgroupReduceLowering.cpp
+++ b/mlir/lib/Dialect/GPU/Transforms/SubgroupReduceLowering.cpp
@@ -11,12 +11,12 @@
//===----------------------------------------------------------------------===//
#include "mlir/Dialect/AMDGPU/IR/AMDGPUDialect.h"
-#include "mlir/Dialect/AMDGPU/Utils/Chipset.h"
#include "mlir/Dialect/Arith/IR/Arith.h"
#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/GPU/Transforms/Passes.h"
#include "mlir/Dialect/GPU/Utils/GPUUtils.h"
#include "mlir/Dialect/LLVMIR/ROCDLDialect.h"
+#include "mlir/Dialect/LLVMIR/ROCDLTargetInfo.h"
#include "mlir/Dialect/Vector/IR/VectorOps.h"
#include "mlir/IR/BuiltinTypes.h"
#include "mlir/IR/Location.h"
@@ -370,7 +370,8 @@ struct VectorSubgroupReduceToShuffles final
static FailureOr<Value>
createSubgroupDPPReduction(PatternRewriter &rewriter, gpu::SubgroupReduceOp op,
Value input, gpu::AllReduceOperation mode,
- const ClusterInfo &ci, amdgpu::Chipset chipset) {
+ const ClusterInfo &ci,
+ const ROCDL::TargetInfo &target) {
Location loc = op.getLoc();
Value dpp;
Value res = input;
@@ -414,7 +415,7 @@ createSubgroupDPPReduction(PatternRewriter &rewriter, gpu::SubgroupReduceOp op,
gpu::convertReductionKind(mode), res, dpp);
}
if (ci.clusterSize >= 32) {
- if (chipset.majorVersion <= 9) {
+ if (!target.has(llvm::AMDGPU::FEAT_GFX10_INSTS)) {
// Broadcast last value from each row to next row.
// Use row mask to avoid polluting row 0 (and row 2 if wave-64).
dpp = amdgpu::DPPOp::create(rewriter, loc, res.getType(), res, res,
@@ -449,7 +450,7 @@ createSubgroupDPPReduction(PatternRewriter &rewriter, gpu::SubgroupReduceOp op,
/*or_mask=*/31,
/*xor_mask=*/0);
}
- } else if (chipset.majorVersion <= 12) {
+ } else if (!target.has(llvm::AMDGPU::FEAT_GFX13_INSTS)) {
// Use a permute lane to cross rows (row 1 <-> row 0, row 3 <-> row 2).
Value uint32Max = arith::ConstantOp::create(
rewriter, loc, rewriter.getI32Type(), rewriter.getI32IntegerAttr(-1));
@@ -472,7 +473,7 @@ createSubgroupDPPReduction(PatternRewriter &rewriter, gpu::SubgroupReduceOp op,
}
}
if (ci.clusterSize >= 64) {
- if (chipset.majorVersion <= 9) {
+ if (!target.has(llvm::AMDGPU::FEAT_GFX10_INSTS)) {
// Broadcast 31st lane value to rows 2 and 3.
dpp = amdgpu::DPPOp::create(rewriter, loc, res.getType(), res, res,
amdgpu::DPPPerm::row_bcast_31,
@@ -486,7 +487,7 @@ createSubgroupDPPReduction(PatternRewriter &rewriter, gpu::SubgroupReduceOp op,
res =
ROCDL::ReadlaneOp::create(rewriter, loc, res.getType(), res, lane63);
- } else if (chipset.majorVersion <= 12) {
+ } else if (!target.has(llvm::AMDGPU::FEAT_GFX13_INSTS)) {
// Assume reduction across 32 lanes has been done.
// Perform final reduction manually by summing values in lane 0 and
// lane 32.
@@ -516,10 +517,11 @@ createSubgroupDPPReduction(PatternRewriter &rewriter, gpu::SubgroupReduceOp op,
struct ScalarSubgroupReduceToDPP final
: OpRewritePattern<gpu::SubgroupReduceOp> {
ScalarSubgroupReduceToDPP(MLIRContext *ctx, unsigned subgroupSize,
- bool matchClustered, amdgpu::Chipset chipset,
+ bool matchClustered,
+ const ROCDL::TargetInfo &target,
PatternBenefit benefit)
: OpRewritePattern(ctx, benefit), subgroupSize(subgroupSize),
- matchClustered(matchClustered), chipset(chipset) {}
+ matchClustered(matchClustered), target(target) {}
LogicalResult matchAndRewrite(gpu::SubgroupReduceOp op,
PatternRewriter &rewriter) const override {
@@ -545,7 +547,7 @@ struct ScalarSubgroupReduceToDPP final
op, "Value type is not a compatible scalar.");
FailureOr<Value> dpp = createSubgroupDPPReduction(
- rewriter, op, op.getValue(), op.getOp(), *ci, chipset);
+ rewriter, op, op.getValue(), op.getOp(), *ci, target);
if (failed(dpp))
return failure();
@@ -556,7 +558,7 @@ struct ScalarSubgroupReduceToDPP final
private:
unsigned subgroupSize = 0;
bool matchClustered = false;
- amdgpu::Chipset chipset;
+ ROCDL::TargetInfo target;
};
} // namespace
@@ -569,18 +571,18 @@ void mlir::populateGpuBreakDownSubgroupReducePatterns(
}
void mlir::populateGpuLowerSubgroupReduceToDPPPatterns(
- RewritePatternSet &patterns, unsigned subgroupSize, amdgpu::Chipset chipset,
- PatternBenefit benefit) {
+ RewritePatternSet &patterns, unsigned subgroupSize,
+ const ROCDL::TargetInfo &target, PatternBenefit benefit) {
patterns.add<ScalarSubgroupReduceToDPP>(patterns.getContext(), subgroupSize,
- /*matchClustered=*/false, chipset,
+ /*matchClustered=*/false, target,
benefit);
}
void mlir::populateGpuLowerClusteredSubgroupReduceToDPPPatterns(
- RewritePatternSet &patterns, unsigned subgroupSize, amdgpu::Chipset chipset,
- PatternBenefit benefit) {
+ RewritePatternSet &patterns, unsigned subgroupSize,
+ const ROCDL::TargetInfo &target, PatternBenefit benefit) {
patterns.add<ScalarSubgroupReduceToDPP>(patterns.getContext(), subgroupSize,
- /*matchClustered=*/true, chipset,
+ /*matchClustered=*/true, target,
benefit);
}
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/8-bit-floats-ocp-gfx1170.mlir b/mlir/test/Conversion/AMDGPUToROCDL/8-bit-floats-ocp-gfx1170.mlir
deleted file mode 100644
index a98d6d1e4c274..0000000000000
--- a/mlir/test/Conversion/AMDGPUToROCDL/8-bit-floats-ocp-gfx1170.mlir
+++ /dev/null
@@ -1,23 +0,0 @@
-// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=chipset=gfx1170 --split-input-file --verify-diagnostics
-
-func.func @ext_packed_fp8(%v: vector<4xf8E4M3FN>) -> f32 {
- // expected-error at below {{failed to legalize operation 'amdgpu.ext_packed_fp8'}}
- %ret = amdgpu.ext_packed_fp8 %v[0] : vector<4xf8E4M3FN> to f32
- func.return %ret : f32
-}
-
-// -----
-
-func.func @ext_packed_bf8(%v: vector<4xf8E5M2>) -> f32 {
- // expected-error at below {{failed to legalize operation 'amdgpu.ext_packed_fp8'}}
- %ret = amdgpu.ext_packed_fp8 %v[0] : vector<4xf8E5M2> to f32
- func.return %ret : f32
-}
-
-// -----
-
-func.func @packed_trunc_2xfp8(%v: f32) -> vector<4xf8E4M3FN> {
- // expected-error at below {{failed to legalize operation 'amdgpu.packed_trunc_2xfp8'}}
- %ret = amdgpu.packed_trunc_2xfp8 %v, undef into undef[word 0] : f32 to vector<4xf8E4M3FN>
- func.return %ret : vector<4xf8E4M3FN>
-}
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/8-bit-floats-ocp.mlir b/mlir/test/Conversion/AMDGPUToROCDL/8-bit-floats-ocp.mlir
index 464d47216c81b..39c340553180e 100644
--- a/mlir/test/Conversion/AMDGPUToROCDL/8-bit-floats-ocp.mlir
+++ b/mlir/test/Conversion/AMDGPUToROCDL/8-bit-floats-ocp.mlir
@@ -1,5 +1,6 @@
-// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=chipset=gfx950 | FileCheck %s
-// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=chipset=gfx1200 | FileCheck %s
+// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=arch=amdgpu9.50-amd-amdhsa | FileCheck %s
+// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=arch=amdgpu12.00-amd-amdhsa | FileCheck %s
+// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=arch=amdgpu11.70-amd-amdhsa | FileCheck %s
// CHECK-LABEL: func @ext_scalar
// CHECK: [[V:%.+]] = builtin.unrealized_conversion_cast %{{.+}} : f8E5M2 to i8
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/8-bit-floats.mlir b/mlir/test/Conversion/AMDGPUToROCDL/8-bit-floats.mlir
index 03fcb266a2e87..c0fc9ed9490f7 100644
--- a/mlir/test/Conversion/AMDGPUToROCDL/8-bit-floats.mlir
+++ b/mlir/test/Conversion/AMDGPUToROCDL/8-bit-floats.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=chipset=gfx942 | FileCheck %s
+// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=arch=amdgpu9.42-amd-amdhsa | FileCheck %s
// CHECK-LABEL: func @ext_scalar
// CHECK: [[V:%.+]] = builtin.unrealized_conversion_cast %{{.+}} : f8E5M2FNUZ to i8
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/amdgpu-to-rocdl.mlir b/mlir/test/Conversion/AMDGPUToROCDL/amdgpu-to-rocdl.mlir
index 8086aa788c8ad..ef82534362bc2 100644
--- a/mlir/test/Conversion/AMDGPUToROCDL/amdgpu-to-rocdl.mlir
+++ b/mlir/test/Conversion/AMDGPUToROCDL/amdgpu-to-rocdl.mlir
@@ -1,10 +1,10 @@
-// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=chipset=gfx908 | FileCheck %s --check-prefixes=CHECK,RECORDS32,GFX9,GFX908
-// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=chipset=gfx90a | FileCheck %s --check-prefixes=CHECK,RECORDS32,GFX9,GFX90A
-// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=chipset=gfx942 | FileCheck %s --check-prefixes=CHECK,RECORDS32,GFX9,GFX942
-// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=chipset=gfx1030 | FileCheck %s --check-prefixes=CHECK,RECORDS32,GFX10,RDNA
-// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=chipset=gfx1100 | FileCheck %s --check-prefixes=CHECK,RECORDS32,GFX11,RDNA
-// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=chipset=gfx1201 | FileCheck %s --check-prefixes=CHECK,RECORDS32,GFX12,RDNA
-// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=chipset=gfx1250 | FileCheck %s --check-prefixes=CHECK,RECORDS45,GFX1250
+// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=arch=amdgpu9.08-amd-amdhsa | FileCheck %s --check-prefixes=CHECK,RECORDS32,GFX9,GFX908
+// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=arch=amdgpu9.0a-amd-amdhsa | FileCheck %s --check-prefixes=CHECK,RECORDS32,GFX9,GFX90A
+// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=arch=amdgpu9.42-amd-amdhsa | FileCheck %s --check-prefixes=CHECK,RECORDS32,GFX9,GFX942
+// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=arch=amdgpu10.30-amd-amdhsa | FileCheck %s --check-prefixes=CHECK,RECORDS32,GFX10,RDNA
+// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=arch=amdgpu11.00-amd-amdhsa | FileCheck %s --check-prefixes=CHECK,RECORDS32,GFX11,RDNA
+// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=arch=amdgpu12.01-amd-amdhsa | FileCheck %s --check-prefixes=CHECK,RECORDS32,GFX12,RDNA
+// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=arch=amdgpu12.50-amd-amdhsa | FileCheck %s --check-prefixes=CHECK,RECORDS45,GFX1250
// CHECK: #[[$MMRA_TAG:.+]] = #llvm.mmra_tag<"amdgpu-synchronize-as":"local">
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/dot-gfx11.mlir b/mlir/test/Conversion/AMDGPUToROCDL/dot-gfx11.mlir
index a87227884dc2a..972577028c3f4 100644
--- a/mlir/test/Conversion/AMDGPUToROCDL/dot-gfx11.mlir
+++ b/mlir/test/Conversion/AMDGPUToROCDL/dot-gfx11.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=chipset=gfx1100 | FileCheck %s
+// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=arch=amdgpu11.00-amd-amdhsa | FileCheck %s
// CHECK-LABEL: @dot_fdot2_f16_f16
func.func @dot_fdot2_f16_f16(%a: vector<2xf16>, %b: vector<2xf16>, %c: f16) -> f16 {
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/dot-gfx12.mlir b/mlir/test/Conversion/AMDGPUToROCDL/dot-gfx12.mlir
index 3213b5fa8f5c2..44d36eb225fb3 100644
--- a/mlir/test/Conversion/AMDGPUToROCDL/dot-gfx12.mlir
+++ b/mlir/test/Conversion/AMDGPUToROCDL/dot-gfx12.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=chipset=gfx1200 | FileCheck %s
+// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=arch=amdgpu12.00-amd-amdhsa | FileCheck %s
// CHECK-LABEL: @dot_fp8_fp8
func.func @dot_fp8_fp8(%a: vector<4xf8E4M3FN>, %b: vector<4xf8E4M3FN>, %c: f32) -> f32 {
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/dot-gfx9.mlir b/mlir/test/Conversion/AMDGPUToROCDL/dot-gfx9.mlir
index e13a9976974dc..022f3bc38ffdd 100644
--- a/mlir/test/Conversion/AMDGPUToROCDL/dot-gfx9.mlir
+++ b/mlir/test/Conversion/AMDGPUToROCDL/dot-gfx9.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=chipset=gfx906 | FileCheck %s
+// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=arch=amdgpu9.06-amd-amdhsa | FileCheck %s
// CHECK-LABEL: @dot_fdot2
func.func @dot_fdot2(%a: vector<2xf16>, %b: vector<2xf16>, %c: f32) -> f32 {
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/dot-invalid.mlir b/mlir/test/Conversion/AMDGPUToROCDL/dot-invalid.mlir
index dd26ab7040734..810a2e02f01ed 100644
--- a/mlir/test/Conversion/AMDGPUToROCDL/dot-invalid.mlir
+++ b/mlir/test/Conversion/AMDGPUToROCDL/dot-invalid.mlir
@@ -1,5 +1,5 @@
-// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=chipset=gfx906 --split-input-file -verify-diagnostics
-// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=chipset=gfx942 --split-input-file -verify-diagnostics
+// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=arch=amdgpu9.06-amd-amdhsa --split-input-file -verify-diagnostics
+// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=arch=amdgpu9.42-amd-amdhsa --split-input-file -verify-diagnostics
// fp8 dot4 is only available on gfx12+.
func.func @dot_fp8_requires_gfx12(%a: vector<4xf8E4M3FN>, %b: vector<4xf8E4M3FN>, %c: f32) -> f32 {
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/dpp.mlir b/mlir/test/Conversion/AMDGPUToROCDL/dpp.mlir
index a4c98111c2956..3c5ccdd02b229 100644
--- a/mlir/test/Conversion/AMDGPUToROCDL/dpp.mlir
+++ b/mlir/test/Conversion/AMDGPUToROCDL/dpp.mlir
@@ -1,6 +1,7 @@
-// RUN: mlir-opt -convert-amdgpu-to-rocdl=chipset=gfx908 %s | FileCheck %s
-// RUN: mlir-opt -convert-amdgpu-to-rocdl=chipset=gfx90a %s | FileCheck %s
-// RUN: mlir-opt -convert-amdgpu-to-rocdl=chipset=gfx942 %s | FileCheck %s
+// RUN: mlir-opt -convert-amdgpu-to-rocdl=arch=amdgpu9.08-amd-amdhsa %s | FileCheck %s
+// RUN: mlir-opt -convert-amdgpu-to-rocdl=arch=amdgpu9.0a-amd-amdhsa %s | FileCheck %s
+// RUN: mlir-opt -convert-amdgpu-to-rocdl=arch=amdgpu9.42-amd-amdhsa %s | FileCheck %s
+// RUN: mlir-opt -convert-amdgpu-to-rocdl=arch=amdgpu9.50-amd-amdhsa %s | FileCheck %s
func.func @test_dpp(%arg0: i32, %arg1: i32) -> i32 {
// CHECK-LABEL: func @test_dpp
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/gfx1250.mlir b/mlir/test/Conversion/AMDGPUToROCDL/gfx1250.mlir
index a98b742c3ce39..1439e97d074f8 100644
--- a/mlir/test/Conversion/AMDGPUToROCDL/gfx1250.mlir
+++ b/mlir/test/Conversion/AMDGPUToROCDL/gfx1250.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=chipset=gfx1250 --split-input-file --verify-diagnostics \
+// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=arch=amdgpu12.50-amd-amdhsa --split-input-file --verify-diagnostics \
// RUN: | FileCheck %s
// CHECK-LABEL: @scaled_ext_packed_matrix_fp4
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/global-prefetch.mlir b/mlir/test/Conversion/AMDGPUToROCDL/global-prefetch.mlir
index f71de64cd071f..1e226ef65b3c6 100644
--- a/mlir/test/Conversion/AMDGPUToROCDL/global-prefetch.mlir
+++ b/mlir/test/Conversion/AMDGPUToROCDL/global-prefetch.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=chipset=gfx1250 --split-input-file --verify-diagnostics | FileCheck %s
+// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=arch=amdgpu12.50-amd-amdhsa --split-input-file --verify-diagnostics | FileCheck %s
// CHECK-LABEL: @glb_prefetch0
func.func @glb_prefetch0(%src : memref<64x64xf16, #gpu.address_space<global>>, %i : i64, %j : i64) {
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/global_transpose_load.mlir b/mlir/test/Conversion/AMDGPUToROCDL/global_transpose_load.mlir
index f378d7232d7b3..a57f2ccfabd4e 100644
--- a/mlir/test/Conversion/AMDGPUToROCDL/global_transpose_load.mlir
+++ b/mlir/test/Conversion/AMDGPUToROCDL/global_transpose_load.mlir
@@ -1,6 +1,6 @@
-// RUN: mlir-opt %s --split-input-file --verify-diagnostics -convert-amdgpu-to-rocdl=chipset=gfx1201 | FileCheck %s
-// RUN: mlir-opt %s --split-input-file -convert-amdgpu-to-rocdl=chipset=gfx1250 | FileCheck %s --check-prefixes=CHECK,CHECK-GFX1250
-// RUN: not mlir-opt %s --split-input-file -convert-amdgpu-to-rocdl=chipset=gfx942 2>&1 | FileCheck %s --check-prefix=CHECK-OLD
+// RUN: mlir-opt %s --split-input-file --verify-diagnostics -convert-amdgpu-to-rocdl=arch=amdgpu12.01-amd-amdhsa | FileCheck %s
+// RUN: mlir-opt %s --split-input-file -convert-amdgpu-to-rocdl=arch=amdgpu12.50-amd-amdhsa | FileCheck %s --check-prefixes=CHECK,CHECK-GFX1250
+// RUN: not mlir-opt %s --split-input-file -convert-amdgpu-to-rocdl=arch=amdgpu9.42-amd-amdhsa 2>&1 | FileCheck %s --check-prefix=CHECK-OLD
// CHECK-LABEL: func @global_transpose_load_8xf16
func.func @global_transpose_load_8xf16(%i : index, %j : index,
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/lds-barrier-gfx90c.mlir b/mlir/test/Conversion/AMDGPUToROCDL/lds-barrier-gfx90c.mlir
index 87b93949be51e..c0eddea2de97c 100644
--- a/mlir/test/Conversion/AMDGPUToROCDL/lds-barrier-gfx90c.mlir
+++ b/mlir/test/Conversion/AMDGPUToROCDL/lds-barrier-gfx90c.mlir
@@ -1,11 +1,12 @@
-// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=chipset=gfx90c | FileCheck %s
+// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=arch=amdgpu9.0c-amd-amdhsa | FileCheck %s
// gfx90c needs the inline assembly workaround, just like gfx908.
// CHECK-LABEL: func @lds_barrier
func.func @lds_barrier() {
// CHECK: llvm.fence syncscope("workgroup") release
- // CHECK-NEXT: rocdl.s.barrier
+ // CHECK-NEXT: llvm.inline_asm has_side_effects asm_dialect = att
+ // CHECK-SAME: ";;;WARNING: BREAKS DEBUG WATCHES\0As_barrier"
// CHECK-NEXT: llvm.fence syncscope("workgroup") acquire
amdgpu.lds_barrier
func.return
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/load_lds-gfx950.mlir b/mlir/test/Conversion/AMDGPUToROCDL/load_lds-gfx950.mlir
index ff8f19c33a437..092ba971169f1 100644
--- a/mlir/test/Conversion/AMDGPUToROCDL/load_lds-gfx950.mlir
+++ b/mlir/test/Conversion/AMDGPUToROCDL/load_lds-gfx950.mlir
@@ -1,5 +1,5 @@
-// RUN: not mlir-opt %s --split-input-file -convert-amdgpu-to-rocdl=chipset=gfx942 2>&1 | FileCheck %s --check-prefix=GFX942
-// RUN: mlir-opt %s --split-input-file -convert-amdgpu-to-rocdl=chipset=gfx950 | FileCheck %s --check-prefix=GFX950
+// RUN: not mlir-opt %s --split-input-file -convert-amdgpu-to-rocdl=arch=amdgpu9.42-amd-amdhsa 2>&1 | FileCheck %s --check-prefix=GFX942
+// RUN: mlir-opt %s --split-input-file -convert-amdgpu-to-rocdl=arch=amdgpu9.50-amd-amdhsa | FileCheck %s --check-prefix=GFX950
// GFX950-LABEL: func @fat_buffer_load_to_rocdl_f96
// GFX950-SAME: (%[[ARG0:.*]]: memref<128x72xf32, #amdgpu.address_space<fat_raw_buffer>>)
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/load_lds.mlir b/mlir/test/Conversion/AMDGPUToROCDL/load_lds.mlir
index c2783c216d66d..2d749e37aadba 100644
--- a/mlir/test/Conversion/AMDGPUToROCDL/load_lds.mlir
+++ b/mlir/test/Conversion/AMDGPUToROCDL/load_lds.mlir
@@ -1,5 +1,5 @@
-// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=chipset=gfx942 | FileCheck %s
-// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=chipset=gfx950 | FileCheck %s
+// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=arch=amdgpu9.42-amd-amdhsa | FileCheck %s
+// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=arch=amdgpu9.50-amd-amdhsa | FileCheck %s
// CHECK-LABEL: func @global_load_to_rocdl_f32
// CHECK-SAME: (%[[ARG0:.*]]: memref<128x72xf32, #gpu.address_space<global>>)
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/memory_counter_wait.mlir b/mlir/test/Conversion/AMDGPUToROCDL/memory_counter_wait.mlir
index 537ef59b503a6..0b2e29d678474 100644
--- a/mlir/test/Conversion/AMDGPUToROCDL/memory_counter_wait.mlir
+++ b/mlir/test/Conversion/AMDGPUToROCDL/memory_counter_wait.mlir
@@ -1,7 +1,7 @@
-// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=chipset=gfx942 | FileCheck %s --check-prefixes=CHECK,GFX9
-// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=chipset=gfx1030 | FileCheck %s --check-prefixes=CHECK,GFX10
-// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=chipset=gfx1100 | FileCheck %s --check-prefixes=CHECK,GFX11
-// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=chipset=gfx1201 | FileCheck %s --check-prefixes=CHECK,GFX12
+// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=arch=amdgpu9.42-amd-amdhsa | FileCheck %s --check-prefixes=CHECK,GFX9
+// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=arch=amdgpu10.30-amd-amdhsa | FileCheck %s --check-prefixes=CHECK,GFX10
+// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=arch=amdgpu11.00-amd-amdhsa | FileCheck %s --check-prefixes=CHECK,GFX11
+// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=arch=amdgpu12.01-amd-amdhsa | FileCheck %s --check-prefixes=CHECK,GFX12
// CHECK-LABEL: func @memory_counter_wait
func.func @memory_counter_wait() {
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/memory_counter_wait_tensor.mlir b/mlir/test/Conversion/AMDGPUToROCDL/memory_counter_wait_tensor.mlir
index 5b29e01abebdb..8e250bef47ad0 100644
--- a/mlir/test/Conversion/AMDGPUToROCDL/memory_counter_wait_tensor.mlir
+++ b/mlir/test/Conversion/AMDGPUToROCDL/memory_counter_wait_tensor.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=chipset=gfx1250 | FileCheck %s
+// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=arch=amdgpu12.50-amd-amdhsa | FileCheck %s
// CHECK-LABEL: func @memory_counter_wait_tensor
func.func @memory_counter_wait_tensor() {
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/memory_counter_wait_unsupported.mlir b/mlir/test/Conversion/AMDGPUToROCDL/memory_counter_wait_unsupported.mlir
index 1d2f692bee488..201fa9caf1e14 100644
--- a/mlir/test/Conversion/AMDGPUToROCDL/memory_counter_wait_unsupported.mlir
+++ b/mlir/test/Conversion/AMDGPUToROCDL/memory_counter_wait_unsupported.mlir
@@ -1,6 +1,6 @@
-// RUN: mlir-opt %s --verify-diagnostics --convert-amdgpu-to-rocdl=chipset=gfx942
-// RUN: mlir-opt %s --verify-diagnostics --convert-amdgpu-to-rocdl=chipset=gfx1030
-// RUN: mlir-opt %s --verify-diagnostics --convert-amdgpu-to-rocdl=chipset=gfx1100
+// RUN: mlir-opt %s --verify-diagnostics --convert-amdgpu-to-rocdl=arch=amdgpu9.42-amd-amdhsa
+// RUN: mlir-opt %s --verify-diagnostics --convert-amdgpu-to-rocdl=arch=amdgpu10.30-amd-amdhsa
+// RUN: mlir-opt %s --verify-diagnostics --convert-amdgpu-to-rocdl=arch=amdgpu11.00-amd-amdhsa
func.func @memory_counter_wait_tensor() {
// expected-error @below{{failed to legalize operation 'amdgpu.memory_counter_wait'}}
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/mfma-fp8-invalid.mlir b/mlir/test/Conversion/AMDGPUToROCDL/mfma-fp8-invalid.mlir
new file mode 100644
index 0000000000000..18d4f2663e40c
--- /dev/null
+++ b/mlir/test/Conversion/AMDGPUToROCDL/mfma-fp8-invalid.mlir
@@ -0,0 +1,21 @@
+// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=arch=amdgpu9.08-amd-amdhsa --split-input-file --verify-diagnostics
+// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=arch=amdgpu9.0a-amd-amdhsa --split-input-file --verify-diagnostics
+
+// gfx908 and gfx90a have FeatureMAIInsts but no fp8 conversions at all, so the
+// fp8 MFMAs -- which first appear on gfx942 -- must not be selected for them.
+
+func.func @mfma_bf8(%arg0 : vector<8xf8E5M2FNUZ>, %arg1 : vector<4xf32>) {
+ // expected-error at below {{op no intrinsic matching MFMA size on given chipset}}
+ // expected-error at below {{failed to legalize operation 'amdgpu.mfma'}}
+ amdgpu.mfma 16x16x32 %arg0 * %arg0 + %arg1 : vector<8xf8E5M2FNUZ>, vector<8xf8E5M2FNUZ>, vector<4xf32>
+ func.return
+}
+
+// -----
+
+func.func @mfma_fp8(%arg0 : vector<8xf8E4M3FNUZ>, %arg1 : vector<4xf32>) {
+ // expected-error at below {{op no intrinsic matching MFMA size on given chipset}}
+ // expected-error at below {{failed to legalize operation 'amdgpu.mfma'}}
+ amdgpu.mfma 16x16x32 %arg0 * %arg0 + %arg1 : vector<8xf8E4M3FNUZ>, vector<8xf8E4M3FNUZ>, vector<4xf32>
+ func.return
+}
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/mfma-gfx950.mlir b/mlir/test/Conversion/AMDGPUToROCDL/mfma-gfx950.mlir
index ba5e096a642f6..8e5bd123121be 100644
--- a/mlir/test/Conversion/AMDGPUToROCDL/mfma-gfx950.mlir
+++ b/mlir/test/Conversion/AMDGPUToROCDL/mfma-gfx950.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=chipset=gfx950 -cse | FileCheck %s
+// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=arch=amdgpu9.50-amd-amdhsa -cse | FileCheck %s
func.func @mfma_to_rocdl(%arg0 : vector<8xf16>, %arg1 : vector<16xf32>,
%arg2 : vector<4xf32>, %arg3 : vector<8xbf16>,
%arg4 : vector<16xi8>, %arg5 : vector<16xi32>,
@@ -96,14 +96,3 @@ func.func @scaled_mfma_to_rocdl(%arg0 : vector<16xf32>,
func.return
}
-
-// CHECK-LABEL: func @mfma_reduce_precision_to_rocdl
-func.func @mfma_reduce_precision_to_rocdl(%arg0 : vector<2xf32>,
- %arg1 : vector<16xf32>,
- %arg2 : vector<4xf32>) {
- // CHECK: rocdl.mfma.f32.32x32x4.xf32
- amdgpu.mfma 32x32x4 %arg0 * %arg0 + %arg1 reducePrecision : vector<2xf32>, vector<2xf32>, vector<16xf32>
- // CHECK: rocdl.mfma.f32.16x16x8.xf32
- amdgpu.mfma 16x16x8 %arg0 * %arg0 + %arg2 reducePrecision : vector<2xf32>, vector<2xf32>, vector<4xf32>
- func.return
-}
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/mfma-reduce-precision-invalid.mlir b/mlir/test/Conversion/AMDGPUToROCDL/mfma-reduce-precision-invalid.mlir
new file mode 100644
index 0000000000000..87e6cd86f869c
--- /dev/null
+++ b/mlir/test/Conversion/AMDGPUToROCDL/mfma-reduce-precision-invalid.mlir
@@ -0,0 +1,23 @@
+// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=arch=amdgpu9.50-amd-amdhsa --split-input-file --verify-diagnostics
+
+// The xf32 MFMAs come from FeatureXF32Insts, which only gfx942 has. gfx950
+// compares greater than gfx942 by ISA version, so a version-ordered check let
+// them through here.
+
+func.func @mfma_reduce_precision_32x32x4(%arg0 : vector<2xf32>,
+ %arg1 : vector<16xf32>) {
+ // expected-error at below {{op no intrinsic matching MFMA size on given chipset}}
+ // expected-error at below {{failed to legalize operation 'amdgpu.mfma'}}
+ amdgpu.mfma 32x32x4 %arg0 * %arg0 + %arg1 reducePrecision : vector<2xf32>, vector<2xf32>, vector<16xf32>
+ func.return
+}
+
+// -----
+
+func.func @mfma_reduce_precision_16x16x8(%arg0 : vector<2xf32>,
+ %arg1 : vector<4xf32>) {
+ // expected-error at below {{op no intrinsic matching MFMA size on given chipset}}
+ // expected-error at below {{failed to legalize operation 'amdgpu.mfma'}}
+ amdgpu.mfma 16x16x8 %arg0 * %arg0 + %arg1 reducePrecision : vector<2xf32>, vector<2xf32>, vector<4xf32>
+ func.return
+}
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/mfma.mlir b/mlir/test/Conversion/AMDGPUToROCDL/mfma.mlir
index 464d2d20048a2..2440851fc099f 100644
--- a/mlir/test/Conversion/AMDGPUToROCDL/mfma.mlir
+++ b/mlir/test/Conversion/AMDGPUToROCDL/mfma.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=chipset=gfx942 -cse | FileCheck %s
+// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=arch=amdgpu9.42-amd-amdhsa -cse | FileCheck %s
func.func @mfma_to_rocdl(%arg0 : f32, %arg1 : vector<32xf32>,
%arg2 : vector<16xf32>, %arg3 : vector<4xf32>,
%arg4 : vector<4xf16>, %arg5 : vector<4xi8>,
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/packed-ext.mlir b/mlir/test/Conversion/AMDGPUToROCDL/packed-ext.mlir
index ad2e7684afc4a..3ba186c8099ae 100644
--- a/mlir/test/Conversion/AMDGPUToROCDL/packed-ext.mlir
+++ b/mlir/test/Conversion/AMDGPUToROCDL/packed-ext.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=chipset=gfx950 | FileCheck %s
+// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=arch=amdgpu9.50-amd-amdhsa | FileCheck %s
// CHECK-LABEL: func.func @scaled_ext_full_f8e4m3_f32
// CHECK-DAG: [[CAST:%.+]] = builtin.unrealized_conversion_cast %arg0 : vector<4xf8E4M3FN> to vector<4xi8>
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/packed-trunc-invalid.mlir b/mlir/test/Conversion/AMDGPUToROCDL/packed-trunc-invalid.mlir
index 93bfc60ce8f4b..6837882d7aa9e 100644
--- a/mlir/test/Conversion/AMDGPUToROCDL/packed-trunc-invalid.mlir
+++ b/mlir/test/Conversion/AMDGPUToROCDL/packed-trunc-invalid.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=chipset=gfx942 --split-input-file --verify-diagnostics
+// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=arch=amdgpu9.42-amd-amdhsa --split-input-file --verify-diagnostics
func.func @packed_trunc_ocp_type_requires_ocp_chipset(%arg0: f32) {
// expected-error at below {{'amdgpu.packed_trunc_2xfp8' op no truncation to result type available on given chipset}}
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/packed-trunc.mlir b/mlir/test/Conversion/AMDGPUToROCDL/packed-trunc.mlir
index e9764d34cefaf..e78f6ac8891ee 100644
--- a/mlir/test/Conversion/AMDGPUToROCDL/packed-trunc.mlir
+++ b/mlir/test/Conversion/AMDGPUToROCDL/packed-trunc.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=chipset=gfx950 | FileCheck %s
+// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=arch=amdgpu9.50-amd-amdhsa | FileCheck %s
// CHECK-LABEL: func.func @packed_scaled_trunc_f8e4m3_f32
// CHECK-DAG: [[ZERO:%.+]] = llvm.mlir.zero : vector<2xi16>
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/permlane-gfx1200-invalid.mlir b/mlir/test/Conversion/AMDGPUToROCDL/permlane-gfx1200-invalid.mlir
new file mode 100644
index 0000000000000..c9745d7fbe08d
--- /dev/null
+++ b/mlir/test/Conversion/AMDGPUToROCDL/permlane-gfx1200-invalid.mlir
@@ -0,0 +1,21 @@
+// RUN: mlir-opt --convert-amdgpu-to-rocdl=arch=amdgpu12.00-amd-amdhsa --split-input-file --verify-diagnostics %s
+
+// gfx1200 has neither FeaturePermlane16Swap nor FeaturePermlane32Swap, but
+// compares greater than gfx950 by ISA version, so a version-ordered check
+// accepted both widths.
+
+func.func @permlane16(%arg0 : i32) -> i32 {
+ // expected-error at below {{op permlane_swap of row length 16 is not supported}}
+ // expected-error at below {{failed to legalize operation 'amdgpu.permlane_swap'}}
+ %0 = amdgpu.permlane_swap %arg0 16 : i32
+ return %0 : i32
+}
+
+// -----
+
+func.func @permlane32(%arg0 : i32) -> i32 {
+ // expected-error at below {{op permlane_swap of row length 32 is not supported}}
+ // expected-error at below {{failed to legalize operation 'amdgpu.permlane_swap'}}
+ %0 = amdgpu.permlane_swap %arg0 32 : i32
+ return %0 : i32
+}
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/permlane-gfx1250-invalid.mlir b/mlir/test/Conversion/AMDGPUToROCDL/permlane-gfx1250-invalid.mlir
new file mode 100644
index 0000000000000..5280aaa90305a
--- /dev/null
+++ b/mlir/test/Conversion/AMDGPUToROCDL/permlane-gfx1250-invalid.mlir
@@ -0,0 +1,11 @@
+// RUN: mlir-opt --convert-amdgpu-to-rocdl=arch=amdgpu12.50-amd-amdhsa --split-input-file --verify-diagnostics %s
+
+// gfx1250 has FeaturePermlane16Swap but not FeaturePermlane32Swap; the 16-wide
+// form is covered as a positive case in permlane.mlir.
+
+func.func @permlane32(%arg0 : i32) -> i32 {
+ // expected-error at below {{op permlane_swap of row length 32 is not supported}}
+ // expected-error at below {{failed to legalize operation 'amdgpu.permlane_swap'}}
+ %0 = amdgpu.permlane_swap %arg0 32 : i32
+ return %0 : i32
+}
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/permlane-gfx1250.mlir b/mlir/test/Conversion/AMDGPUToROCDL/permlane-gfx1250.mlir
new file mode 100644
index 0000000000000..a6b76ce39498b
--- /dev/null
+++ b/mlir/test/Conversion/AMDGPUToROCDL/permlane-gfx1250.mlir
@@ -0,0 +1,12 @@
+// RUN: mlir-opt --convert-amdgpu-to-rocdl=arch=amdgpu12.50-amd-amdhsa --canonicalize %s | FileCheck %s
+
+// gfx1250 has FeaturePermlane16Swap. It does not have FeaturePermlane32Swap;
+// see permlane-gfx1250-invalid.mlir.
+
+// CHECK-LABEL: func @permlane16_i32
+// CHECK-SAME: (%[[ARG0:.*]]: i32)
+func.func @permlane16_i32(%arg0 : i32) -> i32 {
+// CHECK: %[[PERM:.*]] = rocdl.permlane16.swap %[[ARG0]], %[[ARG0]], false, false : (i32, i32) -> <(i32, i32)>
+ %0 = amdgpu.permlane_swap %arg0 16 : i32
+ return %0 : i32
+}
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/permlane-var.mlir b/mlir/test/Conversion/AMDGPUToROCDL/permlane-var.mlir
index e6e5a3061be5b..3ed10ff574bfa 100644
--- a/mlir/test/Conversion/AMDGPUToROCDL/permlane-var.mlir
+++ b/mlir/test/Conversion/AMDGPUToROCDL/permlane-var.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt --convert-amdgpu-to-rocdl=chipset=gfx1200 --canonicalize %s | FileCheck %s
+// RUN: mlir-opt --convert-amdgpu-to-rocdl=arch=amdgpu12.00-amd-amdhsa --canonicalize %s | FileCheck %s
// CHECK-LABEL: func @test_permlane_var_i32
// CHECK-SAME: (%[[SRC:.*]]: i32, %[[SEL:.*]]: i32)
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/permlane.mlir b/mlir/test/Conversion/AMDGPUToROCDL/permlane.mlir
index 9fc49dfeeb9fb..208d82b0d068c 100755
--- a/mlir/test/Conversion/AMDGPUToROCDL/permlane.mlir
+++ b/mlir/test/Conversion/AMDGPUToROCDL/permlane.mlir
@@ -1,5 +1,4 @@
-// RUN: mlir-opt --convert-amdgpu-to-rocdl=chipset=gfx950 --canonicalize %s | FileCheck %s
-// RUN: mlir-opt --convert-amdgpu-to-rocdl=chipset=gfx1200 --canonicalize %s | FileCheck %s
+// RUN: mlir-opt --convert-amdgpu-to-rocdl=arch=amdgpu9.50-amd-amdhsa --canonicalize %s | FileCheck %s
// CHECK-LABEL: func @test_permlane16_i32
// CHECK-SAME: (%[[ARG0:.*]]: i32)
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/sparse-mfma-gfx950.mlir b/mlir/test/Conversion/AMDGPUToROCDL/sparse-mfma-gfx950.mlir
index abdfba9689c8f..3c3efee6546d8 100644
--- a/mlir/test/Conversion/AMDGPUToROCDL/sparse-mfma-gfx950.mlir
+++ b/mlir/test/Conversion/AMDGPUToROCDL/sparse-mfma-gfx950.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=chipset=gfx950 -cse | FileCheck %s
+// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=arch=amdgpu9.50-amd-amdhsa -cse | FileCheck %s
func.func @sparse_mfma_to_rocdl(%arg0 : vector<8xf16>, %arg1 : vector<16xf16>,
%arg2 : vector<4xf32>, %arg3 : vector<16xf32>,
%arg4 : vector<8xbf16>, %arg5 : vector<16xbf16>,
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/sparse-mfma.mlir b/mlir/test/Conversion/AMDGPUToROCDL/sparse-mfma.mlir
index 304f90351faf8..3f49e9dcc80f2 100644
--- a/mlir/test/Conversion/AMDGPUToROCDL/sparse-mfma.mlir
+++ b/mlir/test/Conversion/AMDGPUToROCDL/sparse-mfma.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=chipset=gfx942 -cse | FileCheck %s
+// RUN: mlir-opt %s -convert-amdgpu-to-rocdl=arch=amdgpu9.42-amd-amdhsa -cse | FileCheck %s
func.func @sparse_mfma_to_rocdl(%arg0 : vector<4xf16>, %arg1 : vector<8xf16>,
%arg2 : vector<4xf32>, %arg3 : vector<16xf32>,
%arg4 : vector<4xbf16>, %arg5 : vector<8xbf16>,
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/swizzle.mlir b/mlir/test/Conversion/AMDGPUToROCDL/swizzle.mlir
index ef439efde1bd0..f04e95c92a078 100644
--- a/mlir/test/Conversion/AMDGPUToROCDL/swizzle.mlir
+++ b/mlir/test/Conversion/AMDGPUToROCDL/swizzle.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt -convert-amdgpu-to-rocdl --canonicalize %s | FileCheck %s
+// RUN: mlir-opt -convert-amdgpu-to-rocdl=arch=amdgpu9.42-amd-amdhsa --canonicalize %s | FileCheck %s
// CHECK-LABEL: func @test_swizzle_i32
// CHECK-SAME: (%[[ARG0:.*]]: i32)
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/swmmac-gfx12.mlir b/mlir/test/Conversion/AMDGPUToROCDL/swmmac-gfx12.mlir
index 61d533b75907d..05bbe874e3faa 100644
--- a/mlir/test/Conversion/AMDGPUToROCDL/swmmac-gfx12.mlir
+++ b/mlir/test/Conversion/AMDGPUToROCDL/swmmac-gfx12.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=chipset=gfx1200 --split-input-file --verify-diagnostics | FileCheck %s
+// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=arch=amdgpu12.00-amd-amdhsa --split-input-file --verify-diagnostics | FileCheck %s
// CHECK-LABEL: @rocdl.swmmac
func.func @rocdl.swmmac(
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/swmmac-gfx1250.mlir b/mlir/test/Conversion/AMDGPUToROCDL/swmmac-gfx1250.mlir
index 155e36c369a88..5aae16be1821f 100644
--- a/mlir/test/Conversion/AMDGPUToROCDL/swmmac-gfx1250.mlir
+++ b/mlir/test/Conversion/AMDGPUToROCDL/swmmac-gfx1250.mlir
@@ -1,5 +1,5 @@
-// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=chipset=gfx1250 --split-input-file --verify-diagnostics | FileCheck %s
+// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=arch=amdgpu12.50-amd-amdhsa --split-input-file --verify-diagnostics | FileCheck %s
// CHECK-LABEL: @rocdl.swmmac
func.func @rocdl.swmmac(
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/transpose_load.mlir b/mlir/test/Conversion/AMDGPUToROCDL/transpose_load.mlir
index dcc6624cdb37b..3f1074ebd6b57 100644
--- a/mlir/test/Conversion/AMDGPUToROCDL/transpose_load.mlir
+++ b/mlir/test/Conversion/AMDGPUToROCDL/transpose_load.mlir
@@ -1,5 +1,5 @@
-// RUN: mlir-opt %s --split-input-file -convert-amdgpu-to-rocdl=chipset=gfx950 | FileCheck %s
-// RUN: not mlir-opt %s --split-input-file -convert-amdgpu-to-rocdl=chipset=gfx942 2>&1 | FileCheck %s --check-prefix=CHECK-OLD
+// RUN: mlir-opt %s --split-input-file -convert-amdgpu-to-rocdl=arch=amdgpu9.50-amd-amdhsa | FileCheck %s
+// RUN: not mlir-opt %s --split-input-file -convert-amdgpu-to-rocdl=arch=amdgpu9.42-amd-amdhsa 2>&1 | FileCheck %s --check-prefix=CHECK-OLD
// CHECK-LABEL: func @transpose_load_to_rocdl_4xf16
func.func @transpose_load_to_rocdl_4xf16(%idx1 : index, %idx2 : index, %wgmem : memref<128x72xf16, 3>) -> vector<4xf16> {
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/transpose_load_gfx1250.mlir b/mlir/test/Conversion/AMDGPUToROCDL/transpose_load_gfx1250.mlir
index 98ce6b7ea3001..79ae69eeb2046 100644
--- a/mlir/test/Conversion/AMDGPUToROCDL/transpose_load_gfx1250.mlir
+++ b/mlir/test/Conversion/AMDGPUToROCDL/transpose_load_gfx1250.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt %s --split-input-file -convert-amdgpu-to-rocdl=chipset=gfx1250 | FileCheck %s
+// RUN: mlir-opt %s --split-input-file -convert-amdgpu-to-rocdl=arch=amdgpu12.50-amd-amdhsa | FileCheck %s
// CHECK-LABEL: func @transpose_load_to_rocdl_8xf16
func.func @transpose_load_to_rocdl_8xf16(%idx1 : index, %idx2 : index, %wgmem : memref<128x72xf16, 3>) -> vector<8xf16> {
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/transpose_load_gfx1250_invalid.mlir b/mlir/test/Conversion/AMDGPUToROCDL/transpose_load_gfx1250_invalid.mlir
index 61acdfe245c7c..7cec509a5aef8 100644
--- a/mlir/test/Conversion/AMDGPUToROCDL/transpose_load_gfx1250_invalid.mlir
+++ b/mlir/test/Conversion/AMDGPUToROCDL/transpose_load_gfx1250_invalid.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt %s --split-input-file --verify-diagnostics -convert-amdgpu-to-rocdl=chipset=gfx1250
+// RUN: mlir-opt %s --split-input-file --verify-diagnostics -convert-amdgpu-to-rocdl=arch=amdgpu12.50-amd-amdhsa
func.func @transpose_load_to_rocdl_4xf16(%idx1 : index, %idx2 : index, %wgmem : memref<128x72xf16, 3>) -> vector<4xf16> {
// expected-error at +2 {{'amdgpu.transpose_load' op 16-bit transpose_load requires 8 elements on gfx1250+}}
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/transpose_load_gfx950_invalid.mlir b/mlir/test/Conversion/AMDGPUToROCDL/transpose_load_gfx950_invalid.mlir
index 682f989ede83b..ad685f5807c04 100644
--- a/mlir/test/Conversion/AMDGPUToROCDL/transpose_load_gfx950_invalid.mlir
+++ b/mlir/test/Conversion/AMDGPUToROCDL/transpose_load_gfx950_invalid.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt %s --split-input-file --verify-diagnostics -convert-amdgpu-to-rocdl=chipset=gfx950
+// RUN: mlir-opt %s --split-input-file --verify-diagnostics -convert-amdgpu-to-rocdl=arch=amdgpu9.50-amd-amdhsa
func.func @transpose_load_to_rocdl_8xf16(%idx1 : index, %idx2 : index, %wgmem : memref<128x72xf16, 3>) -> vector<8xf16> {
// expected-error at +2 {{'amdgpu.transpose_load' op 16-bit transpose_load requires 4 elements on gfx950}}
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/transpose_load_reject.mlir b/mlir/test/Conversion/AMDGPUToROCDL/transpose_load_reject.mlir
index a41051c904ed8..08698ca12788e 100644
--- a/mlir/test/Conversion/AMDGPUToROCDL/transpose_load_reject.mlir
+++ b/mlir/test/Conversion/AMDGPUToROCDL/transpose_load_reject.mlir
@@ -1,4 +1,4 @@
-// RUN: not mlir-opt %s --split-input-file -convert-amdgpu-to-rocdl=chipset=gfx950 2>&1 | FileCheck %s
+// RUN: not mlir-opt %s --split-input-file -convert-amdgpu-to-rocdl=arch=amdgpu9.50-amd-amdhsa 2>&1 | FileCheck %s
// -----
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/wmma-gfx11.mlir b/mlir/test/Conversion/AMDGPUToROCDL/wmma-gfx11.mlir
index 08fd68dfe158d..f25b74d660387 100644
--- a/mlir/test/Conversion/AMDGPUToROCDL/wmma-gfx11.mlir
+++ b/mlir/test/Conversion/AMDGPUToROCDL/wmma-gfx11.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=chipset=gfx1100 | FileCheck %s
+// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=arch=amdgpu11.00-amd-amdhsa | FileCheck %s
// CHECK-LABEL: @wmma_to_rocdl
func.func @wmma_to_rocdl(%arg0 : vector<16xf16>, %arg1 : vector<8xf32>, %arg2 : vector<4xf32>,
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/wmma-gfx12.mlir b/mlir/test/Conversion/AMDGPUToROCDL/wmma-gfx12.mlir
index 1dac83946fc4c..dca4df3351d8a 100644
--- a/mlir/test/Conversion/AMDGPUToROCDL/wmma-gfx12.mlir
+++ b/mlir/test/Conversion/AMDGPUToROCDL/wmma-gfx12.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=chipset=gfx1200 \
+// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=arch=amdgpu12.00-amd-amdhsa \
// RUN: --split-input-file --verify-diagnostics | FileCheck %s
// CHECK-LABEL: @wmma_to_rocdl
diff --git a/mlir/test/Conversion/AMDGPUToROCDL/wmma-gfx1250.mlir b/mlir/test/Conversion/AMDGPUToROCDL/wmma-gfx1250.mlir
index 7f9605ad1a7eb..23be3f33e7aee 100644
--- a/mlir/test/Conversion/AMDGPUToROCDL/wmma-gfx1250.mlir
+++ b/mlir/test/Conversion/AMDGPUToROCDL/wmma-gfx1250.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=chipset=gfx1250 \
+// RUN: mlir-opt %s --convert-amdgpu-to-rocdl=arch=amdgpu12.50-amd-amdhsa \
// RUN: --split-input-file --verify-diagnostics | FileCheck %s
// CHECK-LABEL: @wmma_k4
diff --git a/mlir/test/Conversion/ArithToAMDGPU/16-bit-floats.mlir b/mlir/test/Conversion/ArithToAMDGPU/16-bit-floats.mlir
index 6077ef349408f..273f373a9f614 100644
--- a/mlir/test/Conversion/ArithToAMDGPU/16-bit-floats.mlir
+++ b/mlir/test/Conversion/ArithToAMDGPU/16-bit-floats.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt --split-input-file %s -convert-arith-to-amdgpu="allow-packed-f16-round-to-zero=true" | FileCheck %s
+// RUN: mlir-opt --split-input-file %s -convert-arith-to-amdgpu="arch=amdgpu9.42-amd-amdhsa allow-packed-f16-round-to-zero=true" | FileCheck %s
// CHECK-LABEL: @scalar_trunc
// CHECK-SAME: (%[[value:.*]]: f32)
diff --git a/mlir/test/Conversion/ArithToAMDGPU/8-bit-float-saturation-ocp.mlir b/mlir/test/Conversion/ArithToAMDGPU/8-bit-float-saturation-ocp.mlir
index e3c2ae9515939..e4d4aebf87241 100644
--- a/mlir/test/Conversion/ArithToAMDGPU/8-bit-float-saturation-ocp.mlir
+++ b/mlir/test/Conversion/ArithToAMDGPU/8-bit-float-saturation-ocp.mlir
@@ -1,9 +1,9 @@
// RUN: mlir-opt --split-input-file %s \
-// RUN: --pass-pipeline='builtin.module(func.func(convert-arith-to-amdgpu{chipset=gfx950 saturate-fp8-truncf=true}))' \
+// RUN: --pass-pipeline='builtin.module(func.func(convert-arith-to-amdgpu{arch=amdgpu9.50-amd-amdhsa saturate-fp8-truncf=true}))' \
// RUN: | FileCheck %s
// RUN: mlir-opt --split-input-file %s \
-// RUN: --pass-pipeline='builtin.module(func.func(convert-arith-to-amdgpu{chipset=gfx1200 saturate-fp8-truncf=true}))' \
+// RUN: --pass-pipeline='builtin.module(func.func(convert-arith-to-amdgpu{arch=amdgpu12.00-amd-amdhsa saturate-fp8-truncf=true}))' \
// RUN: | FileCheck %s
// CHECK-LABEL: func.func @scalar_trunc
diff --git a/mlir/test/Conversion/ArithToAMDGPU/8-bit-float-saturation.mlir b/mlir/test/Conversion/ArithToAMDGPU/8-bit-float-saturation.mlir
index b6eabe391c0cd..26a0ed993c2fa 100644
--- a/mlir/test/Conversion/ArithToAMDGPU/8-bit-float-saturation.mlir
+++ b/mlir/test/Conversion/ArithToAMDGPU/8-bit-float-saturation.mlir
@@ -1,5 +1,5 @@
// RUN: mlir-opt --split-input-file %s \
-// RUN: --pass-pipeline='builtin.module(func.func(convert-arith-to-amdgpu{chipset=gfx942 saturate-fp8-truncf=true}))' \
+// RUN: --pass-pipeline='builtin.module(func.func(convert-arith-to-amdgpu{arch=amdgpu9.42-amd-amdhsa saturate-fp8-truncf=true}))' \
// RUN: | FileCheck %s
// CHECK-LABEL: func.func @scalar_trunc
diff --git a/mlir/test/Conversion/ArithToAMDGPU/8-bit-floats-ocp.mlir b/mlir/test/Conversion/ArithToAMDGPU/8-bit-floats-ocp.mlir
index 91a9a57898761..a9e2ae202ce73 100644
--- a/mlir/test/Conversion/ArithToAMDGPU/8-bit-floats-ocp.mlir
+++ b/mlir/test/Conversion/ArithToAMDGPU/8-bit-floats-ocp.mlir
@@ -1,6 +1,6 @@
-// RUN: mlir-opt --split-input-file %s -convert-arith-to-amdgpu="chipset=gfx950" | FileCheck %s
-// RUN: mlir-opt --split-input-file %s -convert-arith-to-amdgpu="chipset=gfx1200" | FileCheck %s
-
+// RUN: mlir-opt --split-input-file %s -convert-arith-to-amdgpu="arch=amdgpu9.50-amd-amdhsa" | FileCheck %s
+// RUN: mlir-opt --split-input-file %s -convert-arith-to-amdgpu="arch=amdgpu12.00-amd-amdhsa" | FileCheck %s
+
// CHECK-LABEL: func.func @scalar_ext
// CHECK-SAME: ([[V:%.+]]: f8E5M2)
// CHECK: [[FLOAT:%.+]] = amdgpu.ext_packed_fp8 [[V]][0] : f8E5M2 to f32
diff --git a/mlir/test/Conversion/ArithToAMDGPU/8-bit-floats.mlir b/mlir/test/Conversion/ArithToAMDGPU/8-bit-floats.mlir
index cf3133cf09add..b8f2686f72140 100644
--- a/mlir/test/Conversion/ArithToAMDGPU/8-bit-floats.mlir
+++ b/mlir/test/Conversion/ArithToAMDGPU/8-bit-floats.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt --split-input-file %s -convert-arith-to-amdgpu="chipset=gfx942" | FileCheck %s
+// RUN: mlir-opt --split-input-file %s -convert-arith-to-amdgpu="arch=amdgpu9.42-amd-amdhsa" | FileCheck %s
// CHECK-LABEL: func.func @scalar_ext
// CHECK-SAME: ([[V:%.+]]: f8E5M2FNUZ)
diff --git a/mlir/test/Conversion/ArithToAMDGPU/deprecated-chipset-alias.mlir b/mlir/test/Conversion/ArithToAMDGPU/deprecated-chipset-alias.mlir
new file mode 100644
index 0000000000000..cb72ea0525e74
--- /dev/null
+++ b/mlir/test/Conversion/ArithToAMDGPU/deprecated-chipset-alias.mlir
@@ -0,0 +1,27 @@
+// RUN: mlir-opt %s -convert-arith-to-amdgpu='arch=gfx1030 chipset=gfx942' \
+// RUN: | FileCheck %s --check-prefix=ARCH-WINS
+// RUN: mlir-opt %s -convert-arith-to-amdgpu='chipset=gfx942' \
+// RUN: | FileCheck %s --check-prefix=ALIAS
+
+// Errors name whichever option supplied the target, so a stale `chipset` is
+// reported by its own value rather than by `arch`'s unusable default.
+// RUN: not mlir-opt %s -convert-arith-to-amdgpu='chipset=gfx999' 2>&1 \
+// RUN: | FileCheck %s --check-prefix=BAD-ALIAS
+// BAD-ALIAS: 'gfx999' is not a valid AMDGPU architecture
+
+// With neither given, the unusable default is still what gets reported.
+// RUN: not mlir-opt %s -convert-arith-to-amdgpu 2>&1 \
+// RUN: | FileCheck %s --check-prefix=NEITHER
+// NEITHER: 'invalid' is not a valid AMDGPU architecture
+
+// ARCH-WINS-LABEL: func @truncf_to_fp8
+// ARCH-WINS: arith.truncf
+// ARCH-WINS-NOT: amdgpu.packed_trunc_2xfp8
+
+// ALIAS-LABEL: func @truncf_to_fp8
+// ALIAS: amdgpu.packed_trunc_2xfp8
+// ALIAS-NOT: arith.truncf
+func.func @truncf_to_fp8(%x: f32) -> f8E4M3FNUZ {
+ %r = arith.truncf %x : f32 to f8E4M3FNUZ
+ return %r : f8E4M3FNUZ
+}
diff --git a/mlir/test/Conversion/ArithToAMDGPU/scaling-extf.mlir b/mlir/test/Conversion/ArithToAMDGPU/scaling-extf.mlir
index fe5b0520e37c1..e17b5f2cf16c7 100644
--- a/mlir/test/Conversion/ArithToAMDGPU/scaling-extf.mlir
+++ b/mlir/test/Conversion/ArithToAMDGPU/scaling-extf.mlir
@@ -1,5 +1,5 @@
-// RUN: mlir-opt --split-input-file %s -convert-arith-to-amdgpu="chipset=gfx950" | FileCheck %s
-// RUN: mlir-opt --split-input-file %s -convert-arith-to-amdgpu="chipset=gfx1100" | FileCheck %s --check-prefix=CHECK-GFX1100
+// RUN: mlir-opt --split-input-file %s -convert-arith-to-amdgpu="arch=amdgpu9.50-amd-amdhsa" | FileCheck %s
+// RUN: mlir-opt --split-input-file %s -convert-arith-to-amdgpu="arch=amdgpu11.00-amd-amdhsa" | FileCheck %s --check-prefix=CHECK-GFX1100
// CHECK-LABEL: @conversion_f8_f32_fallback
// CHECK: %[[CST:.+]] = arith.constant dense<0.000000e+00> : vector<2x2xf32>
diff --git a/mlir/test/Conversion/ArithToAMDGPU/scaling-truncf-tensor.mlir b/mlir/test/Conversion/ArithToAMDGPU/scaling-truncf-tensor.mlir
index d22f35a6d07f1..4d94f2a5efb74 100644
--- a/mlir/test/Conversion/ArithToAMDGPU/scaling-truncf-tensor.mlir
+++ b/mlir/test/Conversion/ArithToAMDGPU/scaling-truncf-tensor.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt %s -convert-arith-to-amdgpu=chipset=gfx950 | FileCheck %s
+// RUN: mlir-opt %s -convert-arith-to-amdgpu=arch=amdgpu9.50-amd-amdhsa | FileCheck %s
// CHECK-LABEL: func.func @m0
// CHECK: arith.scaling_truncf
diff --git a/mlir/test/Conversion/ArithToAMDGPU/scaling-truncf.mlir b/mlir/test/Conversion/ArithToAMDGPU/scaling-truncf.mlir
index 4c70768037815..cd26e12be4939 100644
--- a/mlir/test/Conversion/ArithToAMDGPU/scaling-truncf.mlir
+++ b/mlir/test/Conversion/ArithToAMDGPU/scaling-truncf.mlir
@@ -1,5 +1,5 @@
-// RUN: mlir-opt --split-input-file %s -convert-arith-to-amdgpu="chipset=gfx950" | FileCheck %s
-// RUN: mlir-opt --split-input-file %s -convert-arith-to-amdgpu="chipset=gfx1100" | FileCheck %s --check-prefix=CHECK-GFX1100
+// RUN: mlir-opt --split-input-file %s -convert-arith-to-amdgpu="arch=amdgpu9.50-amd-amdhsa" | FileCheck %s
+// RUN: mlir-opt --split-input-file %s -convert-arith-to-amdgpu="arch=amdgpu11.00-amd-amdhsa" | FileCheck %s --check-prefix=CHECK-GFX1100
// CHECK-LABEL: @conversion_f8_fallback
// CHECK-DAG: %[[CST:.+]] = arith.constant dense<0.000000e+00> : vector<2x2xf8E5M2>
diff --git a/mlir/test/Conversion/GPUCommon/lower-global-id.mlir b/mlir/test/Conversion/GPUCommon/lower-global-id.mlir
index 94b9f90052769..18baa4e013f34 100644
--- a/mlir/test/Conversion/GPUCommon/lower-global-id.mlir
+++ b/mlir/test/Conversion/GPUCommon/lower-global-id.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt %s -split-input-file -convert-gpu-to-rocdl | FileCheck %s --check-prefixes=ROCDL
+// RUN: mlir-opt %s -split-input-file -convert-gpu-to-rocdl=arch=amdgpu9.42-amd-amdhsa | FileCheck %s --check-prefixes=ROCDL
// RUN: mlir-opt %s -split-input-file -convert-gpu-to-nvvm | FileCheck %s --check-prefixes=NVVM
gpu.module @kernel {
diff --git a/mlir/test/Conversion/GPUCommon/lower-memory-space-attrs.mlir b/mlir/test/Conversion/GPUCommon/lower-memory-space-attrs.mlir
index 3b794ab717c03..9050122bb7178 100644
--- a/mlir/test/Conversion/GPUCommon/lower-memory-space-attrs.mlir
+++ b/mlir/test/Conversion/GPUCommon/lower-memory-space-attrs.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt %s -split-input-file -convert-gpu-to-rocdl | FileCheck %s --check-prefixes=CHECK,ROCDL
+// RUN: mlir-opt %s -split-input-file -convert-gpu-to-rocdl=arch=amdgpu9.42-amd-amdhsa | FileCheck %s --check-prefixes=CHECK,ROCDL
// RUN: mlir-opt %s -split-input-file -convert-gpu-to-nvvm | FileCheck %s --check-prefixes=CHECK,NVVM
gpu.module @kernel {
diff --git a/mlir/test/Conversion/GPUCommon/memory-attrbution.mlir b/mlir/test/Conversion/GPUCommon/memory-attrbution.mlir
index 38e73ea9179ac..fb3dfd8147436 100644
--- a/mlir/test/Conversion/GPUCommon/memory-attrbution.mlir
+++ b/mlir/test/Conversion/GPUCommon/memory-attrbution.mlir
@@ -1,5 +1,5 @@
// RUN: mlir-opt -allow-unregistered-dialect --convert-gpu-to-nvvm --split-input-file %s | FileCheck --check-prefix=NVVM %s
-// RUN: mlir-opt -allow-unregistered-dialect --convert-gpu-to-rocdl --split-input-file %s | FileCheck --check-prefix=ROCDL %s
+// RUN: mlir-opt -allow-unregistered-dialect --convert-gpu-to-rocdl=arch=amdgpu9.42-amd-amdhsa --split-input-file %s | FileCheck --check-prefix=ROCDL %s
gpu.module @kernel {
// NVVM-LABEL: llvm.func @private
diff --git a/mlir/test/Conversion/GPUCommon/memref-arg-attrs.mlir b/mlir/test/Conversion/GPUCommon/memref-arg-attrs.mlir
index e7c742067b4eb..3646c43313e1f 100644
--- a/mlir/test/Conversion/GPUCommon/memref-arg-attrs.mlir
+++ b/mlir/test/Conversion/GPUCommon/memref-arg-attrs.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt %s -split-input-file -convert-gpu-to-rocdl='use-bare-ptr-memref-call-conv=0' | FileCheck %s --check-prefixes=CHECK,ROCDL
+// RUN: mlir-opt %s -split-input-file -convert-gpu-to-rocdl='arch=amdgpu9.42-amd-amdhsa use-bare-ptr-memref-call-conv=0' | FileCheck %s --check-prefixes=CHECK,ROCDL
// RUN: mlir-opt %s -split-input-file -convert-gpu-to-nvvm='use-bare-ptr-memref-call-conv=0' | FileCheck %s --check-prefixes=CHECK,NVVM
gpu.module @kernel {
diff --git a/mlir/test/Conversion/GPUCommon/memref-arg-noalias-attrs.mlir b/mlir/test/Conversion/GPUCommon/memref-arg-noalias-attrs.mlir
index 33cdc3348e513..0c097fdaa2961 100644
--- a/mlir/test/Conversion/GPUCommon/memref-arg-noalias-attrs.mlir
+++ b/mlir/test/Conversion/GPUCommon/memref-arg-noalias-attrs.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt %s -split-input-file -convert-gpu-to-rocdl='use-bare-ptr-memref-call-conv=1' | FileCheck %s --check-prefixes=CHECK,ROCDL
+// RUN: mlir-opt %s -split-input-file -convert-gpu-to-rocdl='arch=amdgpu9.42-amd-amdhsa use-bare-ptr-memref-call-conv=1' | FileCheck %s --check-prefixes=CHECK,ROCDL
// RUN: mlir-opt %s -split-input-file -convert-gpu-to-nvvm='use-bare-ptr-memref-call-conv=1' | FileCheck %s --check-prefixes=CHECK,NVVM
gpu.module @kernel {
diff --git a/mlir/test/Conversion/GPUCommon/memref-arg-noalias-warning.mlir b/mlir/test/Conversion/GPUCommon/memref-arg-noalias-warning.mlir
index 793df7380d78b..233f21a28b59a 100644
--- a/mlir/test/Conversion/GPUCommon/memref-arg-noalias-warning.mlir
+++ b/mlir/test/Conversion/GPUCommon/memref-arg-noalias-warning.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt %s -split-input-file -convert-gpu-to-rocdl='use-bare-ptr-memref-call-conv=0' -verify-diagnostics
+// RUN: mlir-opt %s -split-input-file -convert-gpu-to-rocdl='arch=amdgpu9.42-amd-amdhsa use-bare-ptr-memref-call-conv=0' -verify-diagnostics
gpu.module @kernel {
// expected-warning @+1 {{Cannot copy noalias with non-bare pointers.}}
diff --git a/mlir/test/Conversion/GPUToROCDL/constant-address-space.mlir b/mlir/test/Conversion/GPUToROCDL/constant-address-space.mlir
index 738aece1769da..9c5ab4bf2edaa 100644
--- a/mlir/test/Conversion/GPUToROCDL/constant-address-space.mlir
+++ b/mlir/test/Conversion/GPUToROCDL/constant-address-space.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt -convert-gpu-to-rocdl %s | FileCheck %s
+// RUN: mlir-opt -convert-gpu-to-rocdl=arch=amdgpu9.42-amd-amdhsa %s | FileCheck %s
module attributes {gpu.container_module} {
gpu.module @kernel_module {
diff --git a/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-barrier.mlir b/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-barrier.mlir
index 618d1889b8478..8da4709b672cd 100644
--- a/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-barrier.mlir
+++ b/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-barrier.mlir
@@ -1,5 +1,5 @@
-// RUN: mlir-opt %s -convert-gpu-to-rocdl='chipset=gfx950' --mlir-print-local-scope | FileCheck %s --check-prefixes=CHECK,GFX9
-// RUN: mlir-opt %s -convert-gpu-to-rocdl='chipset=gfx1201' --mlir-print-local-scope | FileCheck %s --check-prefixes=CHECK,GFX12
+// RUN: mlir-opt %s -convert-gpu-to-rocdl='arch=amdgpu9.50-amd-amdhsa' --mlir-print-local-scope | FileCheck %s --check-prefixes=CHECK,GFX9
+// RUN: mlir-opt %s -convert-gpu-to-rocdl='arch=amdgpu12.01-amd-amdhsa' --mlir-print-local-scope | FileCheck %s --check-prefixes=CHECK,GFX12
gpu.module @test_module {
// CHECK-LABEL: func @barrier_default()
diff --git a/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-barriers-gfx12.mlir b/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-barriers-gfx12.mlir
index 402dcae5e9832..ab4914f331d49 100644
--- a/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-barriers-gfx12.mlir
+++ b/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-barriers-gfx12.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt %s -convert-gpu-to-rocdl='chipset=gfx1250' --mlir-print-local-scope | FileCheck %s
+// RUN: mlir-opt %s -convert-gpu-to-rocdl='arch=amdgpu12.50-amd-amdhsa' --mlir-print-local-scope | FileCheck %s
gpu.module @test_module {
diff --git a/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-hip.mlir b/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-hip.mlir
index 32da31202b688..d2317214d809a 100644
--- a/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-hip.mlir
+++ b/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-hip.mlir
@@ -1,4 +1,6 @@
-// RUN: mlir-opt %s -convert-gpu-to-rocdl='runtime=HIP' -split-input-file | FileCheck %s
+// RUN: mlir-opt %s -convert-gpu-to-rocdl='arch=amdgpu9.42-amd-amdhsa runtime=HIP' -split-input-file | FileCheck %s
+// Ensure old-style options work
+// RUN: mlir-opt %s -convert-gpu-to-rocdl='chipset=gfx942 runtime=HIP' -split-input-file | FileCheck %s
// CHECK-LABEL: gpu.module @test_module
gpu.module @test_module {
diff --git a/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-invalid-ballot.mlir b/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-invalid-ballot.mlir
index a94ab3b5bb780..ad199e413b9b9 100644
--- a/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-invalid-ballot.mlir
+++ b/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-invalid-ballot.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt %s -convert-gpu-to-rocdl='chipset=gfx950' -split-input-file -verify-diagnostics
+// RUN: mlir-opt %s -convert-gpu-to-rocdl='arch=amdgpu9.50-amd-amdhsa' -split-input-file -verify-diagnostics
// -----
diff --git a/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-invalid-dialect.mlir b/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-invalid-dialect.mlir
index 117f7692669de..d917746b12c2d 100644
--- a/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-invalid-dialect.mlir
+++ b/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-invalid-dialect.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt %s -convert-gpu-to-rocdl='allowed-dialects=test' -verify-diagnostics
+// RUN: mlir-opt %s -convert-gpu-to-rocdl='arch=amdgpu9.42-amd-amdhsa allowed-dialects=test' -verify-diagnostics
// expected-error @+1 {{dialect does not implement ConvertToLLVMPatternInterface: test}}
gpu.module @test_module_1 {
diff --git a/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-invalid-named-barrier.mlir b/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-invalid-named-barrier.mlir
index 3f39f4abcf396..9a5750f0660ea 100644
--- a/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-invalid-named-barrier.mlir
+++ b/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-invalid-named-barrier.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt %s -convert-gpu-to-rocdl='chipset=gfx1100' -split-input-file -verify-diagnostics
+// RUN: mlir-opt %s -convert-gpu-to-rocdl='arch=amdgpu11.00-amd-amdhsa' -split-input-file -verify-diagnostics
gpu.module @test_module {
func.func @initialize_named_barrier_pre_gfx12(%count : i32) {
diff --git a/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-named-barrier-non-const.mlir b/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-named-barrier-non-const.mlir
index c9ce2794f1422..cd05533ef9654 100644
--- a/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-named-barrier-non-const.mlir
+++ b/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-named-barrier-non-const.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt %s -convert-gpu-to-rocdl='chipset=gfx1250' -split-input-file -verify-diagnostics
+// RUN: mlir-opt %s -convert-gpu-to-rocdl='arch=amdgpu12.50-amd-amdhsa' -split-input-file -verify-diagnostics
gpu.module @test_module {
func.func @non_constant_member_count(%count : i32) {
diff --git a/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-opencl.mlir b/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-opencl.mlir
index 00d1d7d852680..6f71012a8afda 100644
--- a/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-opencl.mlir
+++ b/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-opencl.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt %s -convert-gpu-to-rocdl='runtime=OpenCL' | FileCheck %s
+// RUN: mlir-opt %s -convert-gpu-to-rocdl='arch=amdgpu9.42-amd-amdhsa runtime=OpenCL' | FileCheck %s
gpu.module @test_module {
// CHECK: llvm.mlir.global internal constant @[[$PRINT_GLOBAL:[A-Za-z0-9_]+]]("Hello: %d\0A\00") {addr_space = 4 : i32}
diff --git a/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-subgroup-id.mlir b/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-subgroup-id.mlir
index 9cab3ff48f5bf..cc474fdaa999f 100644
--- a/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-subgroup-id.mlir
+++ b/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl-subgroup-id.mlir
@@ -1,5 +1,5 @@
-// RUN: mlir-opt %s -convert-gpu-to-rocdl='chipset=gfx942' | FileCheck %s --check-prefixes=CHECK,GFX9
-// RUN: mlir-opt %s -convert-gpu-to-rocdl='chipset=gfx1201' | FileCheck %s --check-prefixes=CHECK,GFX12
+// RUN: mlir-opt %s -convert-gpu-to-rocdl='arch=amdgpu9.42-amd-amdhsa' | FileCheck %s --check-prefixes=CHECK,GFX9
+// RUN: mlir-opt %s -convert-gpu-to-rocdl='arch=amdgpu12.01-amd-amdhsa' | FileCheck %s --check-prefixes=CHECK,GFX12
gpu.module @test_module {
// CHECK-LABEL: func @subgroup_id()
diff --git a/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl.mlir b/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl.mlir
index c01af31e9d4f1..b4051ff32c5f2 100755
--- a/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl.mlir
+++ b/mlir/test/Conversion/GPUToROCDL/gpu-to-rocdl.mlir
@@ -1,6 +1,6 @@
-// RUN: mlir-opt %s -convert-gpu-to-rocdl='chipset=gfx950' -split-input-file | FileCheck %s
-// RUN: mlir-opt %s -convert-gpu-to-rocdl='chipset=gfx950 allowed-dialects=func,arith,math' -split-input-file | FileCheck %s
-// RUN: mlir-opt %s -convert-gpu-to-rocdl='chipset=gfx950 index-bitwidth=32' -split-input-file | FileCheck --check-prefix=CHECK32 %s
+// RUN: mlir-opt %s -convert-gpu-to-rocdl='arch=amdgpu9.50-amd-amdhsa' -split-input-file | FileCheck %s
+// RUN: mlir-opt %s -convert-gpu-to-rocdl='arch=amdgpu9.50-amd-amdhsa allowed-dialects=func,arith,math' -split-input-file | FileCheck %s
+// RUN: mlir-opt %s -convert-gpu-to-rocdl='arch=amdgpu9.50-amd-amdhsa index-bitwidth=32' -split-input-file | FileCheck --check-prefix=CHECK32 %s
// CHECK-LABEL: @test_module
// CHECK-SAME: llvm.data_layout = "e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32-p7:160:256:256:32-p8:128:128:128:48-p9:192:256:256:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-S32-A5-G1-ni:7:8:9"
diff --git a/mlir/test/Conversion/GPUToROCDL/memref.mlir b/mlir/test/Conversion/GPUToROCDL/memref.mlir
index e645481c89230..b6cc3150f003d 100644
--- a/mlir/test/Conversion/GPUToROCDL/memref.mlir
+++ b/mlir/test/Conversion/GPUToROCDL/memref.mlir
@@ -1,6 +1,6 @@
-// RUN: mlir-opt %s -convert-gpu-to-rocdl -split-input-file | FileCheck %s
+// RUN: mlir-opt %s -convert-gpu-to-rocdl=arch=amdgpu9.42-amd-amdhsa -split-input-file | FileCheck %s
// RUN: mlir-opt %s \
-// RUN: -convert-gpu-to-rocdl='use-bare-ptr-memref-call-conv=true' \
+// RUN: -convert-gpu-to-rocdl='arch=amdgpu9.42-amd-amdhsa use-bare-ptr-memref-call-conv=true' \
// RUN: -split-input-file \
// RUN: | FileCheck %s --check-prefix=BARE
diff --git a/mlir/test/Conversion/GPUToROCDL/private-func-no-c-interface.mlir b/mlir/test/Conversion/GPUToROCDL/private-func-no-c-interface.mlir
index a9f778eae2820..b2975cc79404e 100644
--- a/mlir/test/Conversion/GPUToROCDL/private-func-no-c-interface.mlir
+++ b/mlir/test/Conversion/GPUToROCDL/private-func-no-c-interface.mlir
@@ -1,4 +1,4 @@
-// RUN: mlir-opt %s -convert-gpu-to-rocdl -split-input-file | FileCheck %s
+// RUN: mlir-opt %s -convert-gpu-to-rocdl=arch=amdgpu9.42-amd-amdhsa -split-input-file | FileCheck %s
// A private device function cannot be called from outside its module, so it
// must not be given a C interface wrapper; requesting one would only anchor
diff --git a/mlir/test/Conversion/MathToROCDL/math-to-rocdl.mlir b/mlir/test/Conversion/MathToROCDL/math-to-rocdl.mlir
index 455f886839604..b25d30f5d9af8 100644
--- a/mlir/test/Conversion/MathToROCDL/math-to-rocdl.mlir
+++ b/mlir/test/Conversion/MathToROCDL/math-to-rocdl.mlir
@@ -1,5 +1,5 @@
-// RUN: mlir-opt %s -allow-unregistered-dialect -split-input-file -pass-pipeline='builtin.module(convert-math-to-rocdl{chipset=gfx803})' | FileCheck %s --check-prefix=PRE9
-// RUN: mlir-opt %s -allow-unregistered-dialect -split-input-file -pass-pipeline='builtin.module(convert-math-to-rocdl{chipset=gfx942})' | FileCheck %s --check-prefix=POST9
+// RUN: mlir-opt %s -allow-unregistered-dialect -split-input-file -pass-pipeline='builtin.module(convert-math-to-rocdl{arch=amdgpu8.03-amd-amdhsa})' | FileCheck %s --check-prefix=PRE9
+// RUN: mlir-opt %s -allow-unregistered-dialect -split-input-file -pass-pipeline='builtin.module(convert-math-to-rocdl{arch=amdgpu9.42-amd-amdhsa})' | FileCheck %s --check-prefix=POST9
module @test_module {
// CHECK: llvm.func @__ocml_fmod_f16(f16, f16) -> f16
diff --git a/mlir/test/Dialect/AMDGPU/amdgpu-emulate-atomics.mlir b/mlir/test/Dialect/AMDGPU/amdgpu-emulate-atomics.mlir
index fa883bb96c1a0..b4207d82ba296 100644
--- a/mlir/test/Dialect/AMDGPU/amdgpu-emulate-atomics.mlir
+++ b/mlir/test/Dialect/AMDGPU/amdgpu-emulate-atomics.mlir
@@ -1,11 +1,11 @@
-// RUN: mlir-opt -split-input-file -amdgpu-emulate-atomics=chipset=gfx908 %s | FileCheck %s --check-prefixes=CHECK,GFX9CAS,GFX908
-// RUN: mlir-opt -split-input-file -amdgpu-emulate-atomics=chipset=gfx90a %s | FileCheck %s --check-prefixes=CHECK,GFX9CAS,GFX90A
-// RUN: mlir-opt -split-input-file -amdgpu-emulate-atomics=chipset=gfx90c %s | FileCheck %s --check-prefixes=CHECK,GFX9CAS,GFX90C
-// RUN: mlir-opt -split-input-file -amdgpu-emulate-atomics=chipset=gfx1030 %s | FileCheck %s --check-prefixes=CHECK,GFX10
-// RUN: mlir-opt -split-input-file -amdgpu-emulate-atomics=chipset=gfx1100 %s | FileCheck %s --check-prefixes=CHECK,GFX11
-// RUN: mlir-opt -split-input-file -amdgpu-emulate-atomics=chipset=gfx1200 %s | FileCheck %s --check-prefixes=CHECK,GFX12
-// RUN: mlir-opt -split-input-file -amdgpu-emulate-atomics=chipset=gfx942 %s | FileCheck %s --check-prefixes=CHECK,GFX9CAS,GFX942
-// RUN: mlir-opt -split-input-file -amdgpu-emulate-atomics=chipset=gfx950 %s | FileCheck %s --check-prefixes=CHECK,GFX9CAS,GFX950
+// RUN: mlir-opt -split-input-file -amdgpu-emulate-atomics=arch=amdgpu9.08-amd-amdhsa %s | FileCheck %s --check-prefixes=CHECK,GFX9CAS,GFX9NOF64,GFX908
+// RUN: mlir-opt -split-input-file -amdgpu-emulate-atomics=arch=amdgpu9.0a-amd-amdhsa %s | FileCheck %s --check-prefixes=CHECK,GFX9CAS,GFX90A
+// RUN: mlir-opt -split-input-file -amdgpu-emulate-atomics=arch=amdgpu9.0c-amd-amdhsa %s | FileCheck %s --check-prefixes=CHECK,GFX9CAS,GFX9NOF64,GFX90C
+// RUN: mlir-opt -split-input-file -amdgpu-emulate-atomics=arch=amdgpu10.30-amd-amdhsa %s | FileCheck %s --check-prefixes=CHECK,GFX10
+// RUN: mlir-opt -split-input-file -amdgpu-emulate-atomics=arch=amdgpu11.00-amd-amdhsa %s | FileCheck %s --check-prefixes=CHECK,GFX11
+// RUN: mlir-opt -split-input-file -amdgpu-emulate-atomics=arch=amdgpu12.00-amd-amdhsa %s | FileCheck %s --check-prefixes=CHECK,GFX12
+// RUN: mlir-opt -split-input-file -amdgpu-emulate-atomics=arch=amdgpu9.42-amd-amdhsa %s | FileCheck %s --check-prefixes=CHECK,GFX9CAS,GFX942
+// RUN: mlir-opt -split-input-file -amdgpu-emulate-atomics=arch=amdgpu9.50-amd-amdhsa %s | FileCheck %s --check-prefixes=CHECK,GFX9CAS,GFX950
// -----
@@ -46,19 +46,17 @@ func.func @atomic_fmax_f64(%val: f64, %buffer: memref<?xf64>, %idx: i32) {
// GFX12: amdgpu.raw_buffer_atomic_fmax boundsCheck(true) [[val]] -> [[buffer]][[[idx]]]
// GFX942: amdgpu.raw_buffer_atomic_fmax boundsCheck(true) [[val]] -> [[buffer]][[[idx]]]
// GFX950: amdgpu.raw_buffer_atomic_fmax boundsCheck(true) [[val]] -> [[buffer]][[[idx]]]
-// gfx908 has no f64 buffer fmin/fmax, so it is emulated.
-// GFX908: [[ld:%.+]] = amdgpu.raw_buffer_load boundsCheck(true) [[buffer]][[[idx]]]
-// GFX908: cf.br [[loop:\^.+]]([[ld]] : f64)
-// GFX908: [[loop]]([[arg:%.+]]: f64):
-// GFX908: [[operated:%.+]] = arith.maximumf [[val]], [[arg]]
-// GFX908: [[atomicRes:%.+]] = amdgpu.raw_buffer_atomic_cmpswap boundsCheck(true) [[operated]], [[arg]] -> [[buffer]][[[idx]]]
-// GFX908: [[argCast:%.+]] = arith.bitcast [[arg]] : f64 to i64
-// GFX908: [[resCast:%.+]] = arith.bitcast [[atomicRes]] : f64 to i64
-// GFX908: [[test:%.+]] = arith.cmpi eq, [[resCast]], [[argCast]]
-// GFX908: cf.cond_br [[test]], [[post:\^.+]]([[arg]] : f64), [[loop]]([[atomicRes]] : f64)
-// GFX908: [[post]]([[old:%.+]]: f64):
-// gfx90c has none either, but sorts after gfx90a by ISA version.
-// GFX90C: amdgpu.raw_buffer_atomic_fmax boundsCheck(true) [[val]] -> [[buffer]][[[idx]]]
+// Neither gfx908 nor gfx90c has f64 buffer fmin/fmax.
+// GFX9NOF64: [[ld:%.+]] = amdgpu.raw_buffer_load boundsCheck(true) [[buffer]][[[idx]]]
+// GFX9NOF64: cf.br [[loop:\^.+]]([[ld]] : f64)
+// GFX9NOF64: [[loop]]([[arg:%.+]]: f64):
+// GFX9NOF64: [[operated:%.+]] = arith.maximumf [[val]], [[arg]]
+// GFX9NOF64: [[atomicRes:%.+]] = amdgpu.raw_buffer_atomic_cmpswap boundsCheck(true) [[operated]], [[arg]] -> [[buffer]][[[idx]]]
+// GFX9NOF64: [[argCast:%.+]] = arith.bitcast [[arg]] : f64 to i64
+// GFX9NOF64: [[resCast:%.+]] = arith.bitcast [[atomicRes]] : f64 to i64
+// GFX9NOF64: [[test:%.+]] = arith.cmpi eq, [[resCast]], [[argCast]]
+// GFX9NOF64: cf.cond_br [[test]], [[post:\^.+]]([[arg]] : f64), [[loop]]([[atomicRes]] : f64)
+// GFX9NOF64: [[post]]([[old:%.+]]: f64):
// CHECK-NEXT: gpu.printf "End\0A"
gpu.printf "Begin\n"
%old = amdgpu.raw_buffer_atomic_fmax boundsCheck(true) %val -> %buffer[%idx] : f64 -> memref<?xf64>, i32
@@ -77,8 +75,11 @@ func.func @atomic_fadd(%val: f32, %buffer: memref<?xf32>, %idx: i32) {
// GFX12: amdgpu.raw_buffer_atomic_fadd
// GFX942: amdgpu.raw_buffer_atomic_fadd
// GFX950: amdgpu.raw_buffer_atomic_fadd
+// gfx908 only has the no-return form, which suffices here as %old is unused.
// GFX908: amdgpu.raw_buffer_atomic_fadd
-// GFX90C: amdgpu.raw_buffer_atomic_fadd
+// gfx90c has no buffer fadd at all.
+// GFX90C: amdgpu.raw_buffer_load
+// GFX90C: amdgpu.raw_buffer_atomic_cmpswap
%old = amdgpu.raw_buffer_atomic_fadd boundsCheck(true) %val -> %buffer[%idx] : f32 -> memref<?xf32>, i32
func.return
}
@@ -100,8 +101,11 @@ func.func @atomic_fadd_v2f16(%val: vector<2xf16>, %buffer: memref<?xf16>, %idx:
// GFX942: amdgpu.raw_buffer_atomic_fadd
// GFX12: amdgpu.raw_buffer_atomic_fadd
// GFX950: amdgpu.raw_buffer_atomic_fadd
-// GFX908: amdgpu.raw_buffer_atomic_fadd
-// GFX90C: amdgpu.raw_buffer_atomic_fadd
+// Neither gfx908 nor gfx90c has the packed f16 buffer fadd.
+// GFX908: amdgpu.raw_buffer_load
+// GFX908: amdgpu.raw_buffer_atomic_cmpswap
+// GFX90C: amdgpu.raw_buffer_load
+// GFX90C: amdgpu.raw_buffer_atomic_cmpswap
%old = amdgpu.raw_buffer_atomic_fadd boundsCheck(true) %val -> %buffer[%idx] : vector<2xf16> -> memref<?xf16>, i32
func.return
}
@@ -125,3 +129,24 @@ func.func @atomic_fadd_v2bf16(%val: vector<2xbf16>, %buffer: memref<?xbf16>, %id
%old = amdgpu.raw_buffer_atomic_fadd boundsCheck(true) %val -> %buffer[%idx] : vector<2xbf16> -> memref<?xbf16>, i32
func.return
}
+
+// -----
+
+// gfx908 has only the no-return buffer fadd, so a *used* result has to be
+// emulated even though the discarded-result case above lowers natively.
+// CHECK: func @atomic_fadd_used_result
+func.func @atomic_fadd_used_result(%val: f32, %buffer: memref<?xf32>, %idx: i32) -> f32 {
+// GFX908: amdgpu.raw_buffer_load
+// GFX908: amdgpu.raw_buffer_atomic_cmpswap
+// GFX90A: amdgpu.raw_buffer_atomic_fadd
+// GFX942: amdgpu.raw_buffer_atomic_fadd
+// GFX950: amdgpu.raw_buffer_atomic_fadd
+// GFX90C: amdgpu.raw_buffer_load
+// GFX90C: amdgpu.raw_buffer_atomic_cmpswap
+// GFX10: amdgpu.raw_buffer_load
+// GFX10: amdgpu.raw_buffer_atomic_cmpswap
+// GFX11: amdgpu.raw_buffer_atomic_fadd
+// GFX12: amdgpu.raw_buffer_atomic_fadd
+ %old = amdgpu.raw_buffer_atomic_fadd boundsCheck(true) %val -> %buffer[%idx] : f32 -> memref<?xf32>, i32
+ func.return %old : f32
+}
diff --git a/mlir/test/Dialect/GPU/promote-shuffle-amdgpu-invalid.mlir b/mlir/test/Dialect/GPU/promote-shuffle-amdgpu-invalid.mlir
new file mode 100644
index 0000000000000..46a9821cd6fcc
--- /dev/null
+++ b/mlir/test/Dialect/GPU/promote-shuffle-amdgpu-invalid.mlir
@@ -0,0 +1,58 @@
+// RUN: mlir-opt %s --transform-interpreter --split-input-file --verify-diagnostics
+
+module attributes {transform.with_named_sequence} {
+ transform.named_sequence @__transform_main(%module_op: !transform.any_op {transform.readonly}) {
+ %func = transform.structured.match ops{["func.func"]} in %module_op : (!transform.any_op) -> !transform.any_op
+ transform.apply_patterns to %func {
+ // expected-error at below {{'gfx999' is not a valid AMDGPU architecture}}
+ transform.apply_patterns.gpu.gpu_shuffle_to_amdgpu <arch = "gfx999">
+ } : !transform.any_op
+ transform.yield
+ }
+}
+
+// -----
+
+// Only xnack and sramecc are target-ID features, and only on a GPU that
+// supports switching them.
+
+module attributes {transform.with_named_sequence} {
+ transform.named_sequence @__transform_main(%module_op: !transform.any_op {transform.readonly}) {
+ %func = transform.structured.match ops{["func.func"]} in %module_op : (!transform.any_op) -> !transform.any_op
+ transform.apply_patterns to %func {
+ // expected-error at below {{'gfx600:xnack+' is not a valid AMDGPU architecture}}
+ transform.apply_patterns.gpu.gpu_shuffle_to_amdgpu <arch = "gfx600:xnack+">
+ } : !transform.any_op
+ transform.yield
+ }
+}
+
+// -----
+
+// A modifier without a +/- sign is not a target ID.
+
+module attributes {transform.with_named_sequence} {
+ transform.named_sequence @__transform_main(%module_op: !transform.any_op {transform.readonly}) {
+ %func = transform.structured.match ops{["func.func"]} in %module_op : (!transform.any_op) -> !transform.any_op
+ transform.apply_patterns to %func {
+ // expected-error at below {{'gfx908:xnack' is not a valid AMDGPU architecture}}
+ transform.apply_patterns.gpu.gpu_shuffle_to_amdgpu <arch = "gfx908:xnack">
+ } : !transform.any_op
+ transform.yield
+ }
+}
+
+// -----
+
+// Wavefront size is not a target-ID feature, so it cannot ride on `arch`.
+
+module attributes {transform.with_named_sequence} {
+ transform.named_sequence @__transform_main(%module_op: !transform.any_op {transform.readonly}) {
+ %func = transform.structured.match ops{["func.func"]} in %module_op : (!transform.any_op) -> !transform.any_op
+ transform.apply_patterns to %func {
+ // expected-error at below {{'gfx1030:wavefrontsize64+' is not a valid AMDGPU architecture}}
+ transform.apply_patterns.gpu.gpu_shuffle_to_amdgpu <arch = "gfx1030:wavefrontsize64+">
+ } : !transform.any_op
+ transform.yield
+ }
+}
diff --git a/mlir/test/Dialect/GPU/promote-shuffle-amdgpu.mlir b/mlir/test/Dialect/GPU/promote-shuffle-amdgpu.mlir
index 747c997a3b441..d833d8ad58e44 100644
--- a/mlir/test/Dialect/GPU/promote-shuffle-amdgpu.mlir
+++ b/mlir/test/Dialect/GPU/promote-shuffle-amdgpu.mlir
@@ -4,7 +4,7 @@ module attributes {transform.with_named_sequence} {
transform.named_sequence @__transform_main(%module_op: !transform.any_op {transform.readonly}) {
%func = transform.structured.match ops{["func.func"]} in %module_op : (!transform.any_op) -> !transform.any_op
transform.apply_patterns to %func {
- transform.apply_patterns.gpu.gpu_shuffle_to_amdgpu chipset = "gfx950"
+ transform.apply_patterns.gpu.gpu_shuffle_to_amdgpu <arch = "amdgpu9.50-amd-amdhsa">
} : !transform.any_op
transform.yield
}
diff --git a/mlir/test/Dialect/LLVMIR/attach-targets.mlir b/mlir/test/Dialect/LLVMIR/attach-targets.mlir
index eabea4c2bdab5..db507fb15091a 100644
--- a/mlir/test/Dialect/LLVMIR/attach-targets.mlir
+++ b/mlir/test/Dialect/LLVMIR/attach-targets.mlir
@@ -5,7 +5,7 @@
// RUN: | FileCheck %s
// RUN: mlir-opt %s \
// RUN: --nvvm-attach-target='module=options.* O=1 chip=sm_70 fast=true ftz=true' \
-// RUN: --rocdl-attach-target='module=options.* l=file1.bc,file2.bc wave64=false finite-only=true' \
+// RUN: --rocdl-attach-target='module=options.* l=file1.bc,file2.bc wavesize=32 finite-only=true' \
// RUN: --xevm-attach-target='module=options.* O=1 chip=pvc' \
// RUN: | FileCheck %s --check-prefix=CHECK-OPTIONS
diff --git a/mlir/test/Dialect/LLVMIR/rocdl-attach-target-arch.mlir b/mlir/test/Dialect/LLVMIR/rocdl-attach-target-arch.mlir
new file mode 100644
index 0000000000000..b4a1402fdb587
--- /dev/null
+++ b/mlir/test/Dialect/LLVMIR/rocdl-attach-target-arch.mlir
@@ -0,0 +1,105 @@
+// `arch` is an alternative spelling for triple + chip. It is split apart
+// because the attribute feeds the TargetMachine, whose -mcpu only accepts a
+// bare processor name, and the triple is normalized to the new-style spelling
+// that names the GPU's subarch.
+
+// Every spelling of gfx90a lands on the same attribute, whichever triple it
+// arrived with.
+// RUN: mlir-opt %s --rocdl-attach-target='arch=gfx90a' \
+// RUN: | FileCheck %s
+// RUN: mlir-opt %s --rocdl-attach-target='arch=amdgcn-amd-amdhsa--gfx90a' \
+// RUN: | FileCheck %s
+// RUN: mlir-opt %s --rocdl-attach-target='arch=amdgpu9.0a-amd-amdhsa--gfx90a' \
+// RUN: | FileCheck %s
+
+// The legacy split options still attach what they always did, so migrating is
+// opt-in: only `arch` normalizes the triple or consults the target for the
+// wavefront size.
+// RUN: mlir-opt %s \
+// RUN: --rocdl-attach-target='triple=amdgcn-amd-amdhsa chip=gfx90a' \
+// RUN: | FileCheck %s --check-prefix=LEGACY
+
+// An explicit `features` is passed through untouched.
+// RUN: mlir-opt %s --rocdl-attach-target='arch=gfx90a features=+dpp' \
+// RUN: | FileCheck %s --check-prefix=FEATURES
+
+// The processor is the more specific fact, so a family triple normalizes down
+// to the subarch of the GPU it resolved to.
+// RUN: mlir-opt %s --rocdl-attach-target='arch=amdgpu9.4-amd-amdhsa--gfx950' \
+// RUN: | FileCheck %s --check-prefix=FAMILY
+
+// A generic target normalizes to its family subarch.
+// RUN: mlir-opt %s --rocdl-attach-target='arch=gfx9-4-generic' \
+// RUN: | FileCheck %s --check-prefix=GENERIC
+
+// The wavefront size comes from the target, so a wave32-only chip gets
+// no_wave64 rather than the `wave64` option's default.
+// RUN: mlir-opt %s --rocdl-attach-target='arch=gfx1250' \
+// RUN: | FileCheck %s --check-prefix=WAVE32
+
+// A dual-mode chip defaults to wave32 and honours an explicit `wavesize`. The
+// size lands in both the flags (for the device libraries) and the features (for
+// the TargetMachine), which must agree.
+// RUN: mlir-opt %s --rocdl-attach-target='arch=gfx1030' \
+// RUN: | FileCheck %s --check-prefix=WAVE32-RDNA
+// RUN: mlir-opt %s --rocdl-attach-target='arch=gfx1030 wavesize=64' \
+// RUN: | FileCheck %s --check-prefix=WAVE64-RDNA
+
+// Errors have no source location, so they are matched on stderr.
+// RUN: not mlir-opt %s --rocdl-attach-target='arch=gfx999' 2>&1 \
+// RUN: | FileCheck %s --check-prefix=ERR
+// ERR: 'gfx999' is not a valid AMDGPU architecture
+
+// A triple naming no GPU cannot be attached: the attribute requires a chip.
+// RUN: not mlir-opt %s --rocdl-attach-target='arch=amdgcn-amd-amdhsa' 2>&1 \
+// RUN: | FileCheck %s --check-prefix=NOGPU
+// NOGPU: 'amdgcn-amd-amdhsa' names no GPU
+
+// xnack/sramecc are module flags in the backend, not subtarget features, so
+// #rocdl.target has nowhere to carry them and the target ID's modifiers land on
+// the module instead. A modifier the target ID omits stays omitted, since an
+// absent flag means "either" and false would be a different request.
+// RUN: mlir-opt %s --rocdl-attach-target='arch=gfx90a:xnack+' \
+// RUN: | FileCheck %s --check-prefix=XNACK
+// RUN: mlir-opt %s --rocdl-attach-target='arch=gfx90a:sramecc-:xnack-' \
+// RUN: | FileCheck %s --check-prefix=BOTH
+
+module attributes {gpu.container_module} {
+
+// CHECK-LABEL: @rocdl_module
+// CHECK-SAME: [#rocdl.target<triple = "amdgpu9.0a-amd-amdhsa", chip = "gfx90a">]
+// CHECK-NOT: rocdl.xnack
+// CHECK-NOT: rocdl.sramecc
+
+// XNACK-LABEL: @rocdl_module
+// XNACK-SAME: [#rocdl.target<triple = "amdgpu9.0a-amd-amdhsa", chip = "gfx90a">]
+// XNACK-SAME: attributes {rocdl.xnack = true}
+
+// BOTH-LABEL: @rocdl_module
+// BOTH-SAME: [#rocdl.target<triple = "amdgpu9.0a-amd-amdhsa", chip = "gfx90a">]
+// BOTH-SAME: attributes {rocdl.sramecc = false, rocdl.xnack = false}
+
+// LEGACY-LABEL: @rocdl_module
+// LEGACY-SAME: [#rocdl.target<chip = "gfx90a">]
+
+// FEATURES-LABEL: @rocdl_module
+// FEATURES-SAME: [#rocdl.target<triple = "amdgpu9.0a-amd-amdhsa", chip = "gfx90a", features = "+dpp">]
+
+// FAMILY-LABEL: @rocdl_module
+// FAMILY-SAME: [#rocdl.target<triple = "amdgpu9.50-amd-amdhsa", chip = "gfx950">]
+
+// GENERIC-LABEL: @rocdl_module
+// GENERIC-SAME: [#rocdl.target<triple = "amdgpu9.4-amd-amdhsa", chip = "gfx9-4-generic">]
+
+// WAVE32-LABEL: @rocdl_module
+// WAVE32-SAME: [#rocdl.target<triple = "amdgpu12.50-amd-amdhsa", chip = "gfx1250", flags = {no_wave64}>]
+
+// WAVE32-RDNA-LABEL: @rocdl_module
+// WAVE32-RDNA-SAME: [#rocdl.target<triple = "amdgpu10.30-amd-amdhsa", chip = "gfx1030", features = "+wavefrontsize32", flags = {no_wave64}>]
+
+// WAVE64-RDNA-LABEL: @rocdl_module
+// WAVE64-RDNA-SAME: [#rocdl.target<triple = "amdgpu10.30-amd-amdhsa", chip = "gfx1030", features = "+wavefrontsize64">]
+gpu.module @rocdl_module {
+}
+
+}
diff --git a/mlir/test/Integration/GPU/ROCM/gpu-lower-to-rocdl-pipeline.mlir b/mlir/test/Integration/GPU/ROCM/gpu-lower-to-rocdl-pipeline.mlir
index ee289b9cd5549..c9d8561d2e82d 100644
--- a/mlir/test/Integration/GPU/ROCM/gpu-lower-to-rocdl-pipeline.mlir
+++ b/mlir/test/Integration/GPU/ROCM/gpu-lower-to-rocdl-pipeline.mlir
@@ -1,5 +1,5 @@
// RUN: mlir-opt %s \
-// RUN: --gpu-lower-to-rocdl-pipeline="chip=%chip" \
+// RUN: --gpu-lower-to-rocdl-pipeline="arch=%chip" \
// RUN: | mlir-runner \
// RUN: --shared-libs=%mlir_rocm_runtime \
// RUN: --shared-libs=%mlir_runner_utils \
diff --git a/mlir/test/Integration/GPU/ROCM/gpu-to-hsaco.mlir b/mlir/test/Integration/GPU/ROCM/gpu-to-hsaco.mlir
index 5fa27eab3bba4..ea570c2970b71 100644
--- a/mlir/test/Integration/GPU/ROCM/gpu-to-hsaco.mlir
+++ b/mlir/test/Integration/GPU/ROCM/gpu-to-hsaco.mlir
@@ -1,6 +1,6 @@
// RUN: mlir-opt %s \
// RUN: | mlir-opt -gpu-kernel-outlining \
-// RUN: | mlir-opt -pass-pipeline='builtin.module(gpu.module(strip-debuginfo,convert-gpu-to-rocdl),rocdl-attach-target{chip=%chip})' \
+// RUN: | mlir-opt -pass-pipeline='builtin.module(gpu.module(strip-debuginfo,convert-gpu-to-rocdl{arch=%chip}),rocdl-attach-target{arch=%chip})' \
// RUN: | mlir-opt -gpu-to-llvm -reconcile-unrealized-casts -gpu-module-to-binary \
// RUN: | mlir-runner \
// RUN: --shared-libs=%mlir_rocm_runtime \
diff --git a/mlir/test/Integration/GPU/ROCM/printf.mlir b/mlir/test/Integration/GPU/ROCM/printf.mlir
index 8327ec428589d..1c67dae8cfe27 100644
--- a/mlir/test/Integration/GPU/ROCM/printf.mlir
+++ b/mlir/test/Integration/GPU/ROCM/printf.mlir
@@ -1,5 +1,5 @@
// RUN: mlir-opt %s \
-// RUN: | mlir-opt -pass-pipeline='builtin.module(gpu.module(strip-debuginfo,convert-gpu-to-rocdl{index-bitwidth=32 runtime=HIP}),rocdl-attach-target{chip=%chip})' \
+// RUN: | mlir-opt -pass-pipeline='builtin.module(gpu.module(strip-debuginfo,convert-gpu-to-rocdl{arch=%chip index-bitwidth=32 runtime=HIP}),rocdl-attach-target{arch=%chip})' \
// RUN: | mlir-opt -gpu-to-llvm -reconcile-unrealized-casts -gpu-module-to-binary \
// RUN: | mlir-runner \
// RUN: --shared-libs=%mlir_rocm_runtime \
diff --git a/mlir/test/Integration/GPU/ROCM/two-modules.mlir b/mlir/test/Integration/GPU/ROCM/two-modules.mlir
index f3062dbda86c8..688fd2248b835 100644
--- a/mlir/test/Integration/GPU/ROCM/two-modules.mlir
+++ b/mlir/test/Integration/GPU/ROCM/two-modules.mlir
@@ -1,6 +1,6 @@
// RUN: mlir-opt %s \
// RUN: | mlir-opt -gpu-kernel-outlining \
-// RUN: | mlir-opt -pass-pipeline='builtin.module(gpu.module(strip-debuginfo,convert-gpu-to-rocdl),rocdl-attach-target{chip=%chip})' \
+// RUN: | mlir-opt -pass-pipeline='builtin.module(gpu.module(strip-debuginfo,convert-gpu-to-rocdl{arch=%chip}),rocdl-attach-target{arch=%chip})' \
// RUN: | mlir-opt -gpu-to-llvm -reconcile-unrealized-casts -gpu-module-to-binary \
// RUN: | mlir-runner \
// RUN: --shared-libs=%mlir_rocm_runtime \
diff --git a/mlir/test/Integration/GPU/ROCM/vecadd.mlir b/mlir/test/Integration/GPU/ROCM/vecadd.mlir
index c3f8b982a131a..88a969fef8201 100644
--- a/mlir/test/Integration/GPU/ROCM/vecadd.mlir
+++ b/mlir/test/Integration/GPU/ROCM/vecadd.mlir
@@ -1,7 +1,7 @@
// RUN: mlir-opt %s \
// RUN: | mlir-opt -convert-scf-to-cf \
// RUN: | mlir-opt -gpu-kernel-outlining \
-// RUN: | mlir-opt -pass-pipeline='builtin.module(gpu.module(strip-debuginfo,convert-gpu-to-rocdl{use-bare-ptr-memref-call-conv=true}),rocdl-attach-target{chip=%chip})' \
+// RUN: | mlir-opt -pass-pipeline='builtin.module(gpu.module(strip-debuginfo,convert-gpu-to-rocdl{arch=%chip use-bare-ptr-memref-call-conv=true}),rocdl-attach-target{arch=%chip})' \
// RUN: | mlir-opt -gpu-to-llvm=use-bare-pointers-for-kernels=true -reconcile-unrealized-casts -gpu-module-to-binary \
// RUN: | mlir-runner \
// RUN: --shared-libs=%mlir_rocm_runtime \
diff --git a/mlir/test/Integration/GPU/ROCM/vector-transferops.mlir b/mlir/test/Integration/GPU/ROCM/vector-transferops.mlir
index a633edb8377af..6c13ded50307c 100644
--- a/mlir/test/Integration/GPU/ROCM/vector-transferops.mlir
+++ b/mlir/test/Integration/GPU/ROCM/vector-transferops.mlir
@@ -1,7 +1,7 @@
// RUN: mlir-opt %s \
// RUN: | mlir-opt -convert-scf-to-cf \
// RUN: | mlir-opt -gpu-kernel-outlining \
-// RUN: | mlir-opt -pass-pipeline='builtin.module(gpu.module(strip-debuginfo,convert-gpu-to-rocdl{chipset=%chip index-bitwidth=32}),rocdl-attach-target{chip=%chip})' \
+// RUN: | mlir-opt -pass-pipeline='builtin.module(gpu.module(strip-debuginfo,convert-gpu-to-rocdl{arch=%chip index-bitwidth=32}),rocdl-attach-target{arch=%chip})' \
// RUN: | mlir-opt -gpu-to-llvm -reconcile-unrealized-casts -gpu-module-to-binary \
// RUN: | mlir-runner \
// RUN: --shared-libs=%mlir_rocm_runtime \
diff --git a/mlir/test/lib/Dialect/GPU/TestGpuRewrite.cpp b/mlir/test/lib/Dialect/GPU/TestGpuRewrite.cpp
index 72c3952f0f496..4b620496d6b41 100644
--- a/mlir/test/lib/Dialect/GPU/TestGpuRewrite.cpp
+++ b/mlir/test/lib/Dialect/GPU/TestGpuRewrite.cpp
@@ -11,7 +11,6 @@
//===----------------------------------------------------------------------===//
#include "mlir/Dialect/AMDGPU/IR/AMDGPUDialect.h"
-#include "mlir/Dialect/AMDGPU/Utils/Chipset.h"
#include "mlir/Dialect/Arith/IR/Arith.h"
#include "mlir/Dialect/Func/IR/FuncDialect.h"
#include "mlir/Dialect/GPU/Transforms/Passes.h"
@@ -91,12 +90,16 @@ struct TestGpuSubgroupReduceLoweringPass
/*maxShuffleBitwidth=*/32,
PatternBenefit(3));
if (expandToShuffles) {
- auto maybeChipset = amdgpu::Chipset::parse(target);
- if (succeeded(maybeChipset)) {
+ if (!target.empty()) {
+ FailureOr<ROCDL::TargetInfo> targetInfo =
+ ROCDL::TargetInfo::get(target, /*waveSize=*/0,
+ [&] { return getOperation()->emitError(); });
+ if (failed(targetInfo))
+ return signalPassFailure();
populateGpuLowerSubgroupReduceToDPPPatterns(
- patterns, /*subgroupSize=*/64, *maybeChipset, PatternBenefit(2));
+ patterns, /*subgroupSize=*/64, *targetInfo, PatternBenefit(2));
populateGpuLowerClusteredSubgroupReduceToDPPPatterns(
- patterns, /*subgroupSize=*/64, *maybeChipset, PatternBenefit(2));
+ patterns, /*subgroupSize=*/64, *targetInfo, PatternBenefit(2));
}
populateGpuLowerSubgroupReduceToShufflePatterns(
patterns, /*subgroupSize=*/32, /*shuffleBitwidth=*/32);
diff --git a/mlir/unittests/Dialect/AMDGPU/AMDGPUUtilsTest.cpp b/mlir/unittests/Dialect/AMDGPU/AMDGPUUtilsTest.cpp
index 570d56f3c6ff1..33fbc79e77027 100644
--- a/mlir/unittests/Dialect/AMDGPU/AMDGPUUtilsTest.cpp
+++ b/mlir/unittests/Dialect/AMDGPU/AMDGPUUtilsTest.cpp
@@ -9,6 +9,14 @@
#include "mlir/Dialect/AMDGPU/Utils/Chipset.h"
#include "gtest/gtest.h"
+// Chipset is deprecated in favour of ROCDL::TargetInfo, but stays covered for
+// as long as it ships.
+#ifdef __clang__
+#pragma clang diagnostic ignored "-Wdeprecated-declarations"
+#elif defined(__GNUC__)
+#pragma GCC diagnostic ignored "-Wdeprecated-declarations"
+#endif
+
namespace mlir::amdgpu {
namespace {
More information about the llvm-branch-commits
mailing list