[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
Wed Sep 16 11:47:21 PDT 2026
https://github.com/krzysz00 updated https://github.com/llvm/llvm-project/pull/223563
>From 1cc3e3128c97ad8cb496ca63ff8fb56d76399b80 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 16b93c89096709..b4650b82841b3d 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 393658652dbac4..861bbe0cc8f754 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 fd144edf77452a..e6e137759a76a3 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 220da0ad3c08fe..494ee8cfa11abd 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 60f1888569362b..8ba104972abffe 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 6f3e288d454fed..2f80414d70b1f5 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 48e7658568f86f..c8ec368231f835 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 7dd7ac750a9eb2..d9e68a7be4dc7c 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 256065537aeb9b..1a964c27a63335 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 6327415d62769f..c34cc0ee9fdd6f 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 030506bba10768..d56cc1f13afd2e 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 d5c253d6c9c08c..6657ff819dea5c 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 1aaab1ee915235..090bfdba3ce8f5 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 6bbe65731964f1..4cb739816e97df 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 d0714739589f5c..4a5399d7bf4bde 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 c6f81fd533dffd..94f0f82b2c39ed 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 04c3c9e4136e71..b54de7e87adabd 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 29baef635ec802..28fa958181d038 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 ce47da8d9ee973..930235995b4e1a 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 1e5fd09a00a758..edbcaa1f8ca1c0 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 9591b76a5330ce..a7858781a10792 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 01da26f88a84d0..8fb62258c98ea3 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 1f44ffa52e068e..cef6deee5c3548 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 ec1571a56fe4a8..23f1a16ca85aa7 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 a98d6d1e4c2745..00000000000000
--- 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 464d47216c81be..39c340553180e2 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 03fcb266a2e870..c0fc9ed9490f7c 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 8086aa788c8add..ef82534362bc21 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 a87227884dc2a9..972577028c3f4b 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 3213b5fa8f5c2f..44d36eb225fb3f 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 e13a9976974dc6..022f3bc38ffdd4 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 dd26ab7040734c..810a2e02f01eda 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 a4c98111c29561..3c5ccdd02b2298 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 a98b742c3ce39a..1439e97d074f8e 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 f71de64cd071f2..1e226ef65b3c67 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 f378d7232d7b33..a57f2ccfabd4e4 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 87b93949be51e0..c0eddea2de97c2 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 ff8f19c33a437d..092ba971169f10 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 c2783c216d66d9..2d749e37aadba7 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 537ef59b503a6e..0b2e29d678474e 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 5b29e01abebdbc..8e250bef47ad0d 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 1d2f692bee4882..201fa9caf1e14d 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 00000000000000..18d4f2663e40c8
--- /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 ba5e096a642f62..8e5bd123121bec 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 00000000000000..87e6cd86f869cc
--- /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 464d2d20048a27..2440851fc099f7 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 ad2e7684afc4a3..3ba186c8099ae1 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 93bfc60ce8f4be..6837882d7aa9e1 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 e9764d34cefaf8..e78f6ac8891ee5 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 00000000000000..c9745d7fbe08d0
--- /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 00000000000000..5280aaa90305a1
--- /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 00000000000000..a6b76ce39498ba
--- /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 e6e5a3061be5b0..3ed10ff574bfa4 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 9fc49dfeeb9fb5..208d82b0d068c1 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 abdfba9689c8fa..3c3efee6546d8c 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 304f90351faf8b..3f49e9dcc80f22 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 ef439efde1bd02..f04e95c92a0781 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 61d533b75907d4..05bbe874e3faa3 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 155e36c369a88b..5aae16be1821f6 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 dcc6624cdb37b8..3f1074ebd6b571 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 98ce6b7ea3001e..79ae69eeb2046f 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 61acdfe245c7cb..7cec509a5aef89 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 682f989ede83be..ad685f5807c041 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 a41051c904ed8d..08698ca12788e9 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 08fd68dfe158de..f25b74d6603879 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 1dac83946fc4cf..dca4df3351d8aa 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 7f9605ad1a7eb9..23be3f33e7aee7 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 6077ef349408f9..273f373a9f614a 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 e3c2ae95159395..e4d4aebf87241b 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 b6eabe391c0cd7..26a0ed993c2fac 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 91a9a57898761c..a9e2ae202ce735 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 cf3133cf09add9..b8f2686f721409 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 00000000000000..cb72ea0525e741
--- /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 fe5b0520e37c16..e17b5f2cf16c76 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 d22f35a6d07f1d..4d94f2a5efb743 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 4c707680378158..cd26e12be49399 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 94b9f90052769e..18baa4e013f347 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 3b794ab717c03f..9050122bb71785 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 38e73ea9179ac6..fb3dfd81474360 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 e7c742067b4eb5..3646c43313e1f9 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 33cdc3348e5137..0c097fdaa29613 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 793df7380d78bd..233f21a28b59ab 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 738aece1769dab..9c5ab4bf2edaa9 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 618d1889b84785..8da4709b672cd8 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 402dcae5e98323..ab4914f331d495 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 32da31202b6884..d2317214d809a2 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 a94ab3b5bb780b..ad199e413b9b9f 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 117f7692669de8..d917746b12c2dc 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 3f39f4abcf396a..9a5750f0660ea9 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 c9ce2794f1422a..cd05533ef96542 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 00d1d7d8526809..6f71012a8afda5 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 9cab3ff48f5bf0..cc474fdaa999f6 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 c01af31e9d4f15..b4051ff32c5f21 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 e645481c892308..b6cc3150f003d2 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 a9f778eae28209..b2975cc79404e2 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 455f8868396044..b25d30f5d9af87 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 fa883bb96c1a04..b4207d82ba2960 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 00000000000000..46a9821cd6fcc5
--- /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 747c997a3b441b..d833d8ad58e445 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 eabea4c2bdab54..db507fb15091a9 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 00000000000000..b4a1402fdb5878
--- /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 ee289b9cd55491..c9d8561d2e82d6 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 5fa27eab3bba46..ea570c2970b71d 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 8327ec428589dd..1c67dae8cfe27a 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 f3062dbda86c8f..688fd2248b835b 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 c3f8b982a131a1..88a969fef82018 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 a633edb8377af9..6c13ded50307c1 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 72c3952f0f496b..4b620496d6b410 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 570d56f3c6ff13..33fbc79e770277 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