[ROCM] Name gfx1250 CDNA5 and add MI455X target details (#24923)
AMD published the CDNA5 architecture and MI455X specifications a few
weeks ago:
https://www.amd.com/content/dam/amd/en/documents/products/technologies/cdna/amd-cdna5-whitepaper.pdf
Move the existing gfx1250 WGP definition alongside the other CDNA
families and add cdna5 and mi455x target aliases. Populate the MI455X
chip definition with the published WGP count, memory bandwidth, and peak
compute rates. Document the confirmed WGP limits and update the
architecture naming in compiler and runtime documentation.
Normalize SKU and architecture aliases before forming executable target
IDs for the amdgpu device. Without this, mi455x and cdna5 are emitted as
executable formats that the runtime cannot match, even though the
embedded target_info already names gfx1250.
Add coverage for target metadata, Wave32 serialization, and canonical
AMDGPU executable formats with feature suffixes.
Assisted-by: codex (gpt-6-astra)
diff --git a/compiler/plugins/target/ROCM/BUILD.bazel b/compiler/plugins/target/ROCM/BUILD.bazel
index c1190b2..ed8ae58 100644
--- a/compiler/plugins/target/ROCM/BUILD.bazel
+++ b/compiler/plugins/target/ROCM/BUILD.bazel
@@ -51,6 +51,7 @@
"//runtime/src/iree/schemas:hip_executable_def_c_fbs",
"@llvm-project//llvm:AMDGPUCodeGen",
"@llvm-project//llvm:Analysis",
+ "@llvm-project//llvm:BinaryFormat",
"@llvm-project//llvm:BitWriter",
"@llvm-project//llvm:Core",
"@llvm-project//llvm:FrontendOffloading",
@@ -62,6 +63,7 @@
"@llvm-project//llvm:SPIRVCodeGen",
"@llvm-project//llvm:Support",
"@llvm-project//llvm:Target",
+ "@llvm-project//llvm:TargetParser",
"@llvm-project//llvm:TransformUtils",
"@llvm-project//mlir:AMDGPUDialect",
"@llvm-project//mlir:AMDGPUToROCDL",
diff --git a/compiler/plugins/target/ROCM/CMakeLists.txt b/compiler/plugins/target/ROCM/CMakeLists.txt
index 442103a..7a77b04 100644
--- a/compiler/plugins/target/ROCM/CMakeLists.txt
+++ b/compiler/plugins/target/ROCM/CMakeLists.txt
@@ -28,6 +28,7 @@
DEPS
LLVMAMDGPUCodeGen
LLVMAnalysis
+ LLVMBinaryFormat
LLVMBitWriter
LLVMCore
LLVMFrontendOffloading
@@ -38,6 +39,7 @@
LLVMSPIRVCodeGen
LLVMSupport
LLVMTarget
+ LLVMTargetParser
LLVMTransformUtils
LLVMipo
MLIRAMDGPUDialect
diff --git a/compiler/plugins/target/ROCM/ROCMTarget.cpp b/compiler/plugins/target/ROCM/ROCMTarget.cpp
index 2775a1d..1b83dcd 100644
--- a/compiler/plugins/target/ROCM/ROCMTarget.cpp
+++ b/compiler/plugins/target/ROCM/ROCMTarget.cpp
@@ -6,6 +6,7 @@
#include "ROCMTargetUtils.h"
+#include <cstddef>
#include <cstdint>
#include "compiler/plugins/target/ROCM/Dialect/ROCM/IR/ROCMAttrs.h"
@@ -37,6 +38,7 @@
#include "llvm/ADT/StringExtras.h"
#include "llvm/ADT/StringRef.h"
#include "llvm/Analysis/TargetTransformInfo.h"
+#include "llvm/BinaryFormat/ELF.h"
#include "llvm/Bitcode/BitcodeWriter.h"
#include "llvm/Frontend/Offloading/Utility.h"
#include "llvm/IR/Constants.h"
@@ -48,9 +50,11 @@
#include "llvm/Passes/PassBuilder.h"
#include "llvm/Passes/StandardInstrumentations.h"
#include "llvm/Support/CommandLine.h"
+#include "llvm/Support/Endian.h"
#include "llvm/Support/FormatVariadic.h"
#include "llvm/Support/TargetSelect.h"
#include "llvm/Target/TargetMachine.h"
+#include "llvm/TargetParser/AMDGPUTargetParser.h"
#include "llvm/Transforms/Utils/Cloning.h"
#include "mlir/Dialect/AMDGPU/Utils/Chipset.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
@@ -165,9 +169,16 @@
}
}
-static FailureOr<std::string> buildAMDGPUTargetID(Location loc,
- StringRef targetArch,
- StringRef targetFeatures) {
+static llvm::Triple getAMDGPUTargetTriple(StringRef targetArch) {
+ llvm::Triple triple("amdgcn-amd-amdhsa");
+ triple.setArch(llvm::Triple::amdgpu,
+ llvm::AMDGPU::getSubArchFromGPUName(targetArch));
+ return triple;
+}
+
+static FailureOr<llvm::AMDGPU::TargetID>
+parseAMDGPUTargetID(Location loc, StringRef targetArch,
+ StringRef targetFeatures) {
FailureOr<AMDGPUTargetFeatureModes> modes =
parseAMDGPUTargetFeatureModes(loc, targetFeatures);
if (failed(modes)) {
@@ -176,7 +187,12 @@
std::string targetID = targetArch.str();
appendAMDGPUTargetFeatureSuffix(targetID, "sramecc", modes->sramecc);
appendAMDGPUTargetFeatureSuffix(targetID, "xnack", modes->xnack);
- return targetID;
+ auto parsedTargetID = llvm::AMDGPU::TargetID::parse(
+ getAMDGPUTargetTriple(targetArch), targetID);
+ if (!parsedTargetID) {
+ return emitError(loc) << "invalid ROCM target ID '" << targetID << "'";
+ }
+ return *parsedTargetID;
}
struct ROCMOptions {
@@ -310,8 +326,9 @@
return emitError(builder.getUnknownLoc(), "Unknown ROCM target '")
<< target << "'";
}
- if (failed(parseAMDGPUTargetFeatureModes(builder.getUnknownLoc(),
- targetFeatures))) {
+ if (failed(parseAMDGPUTargetID(builder.getUnknownLoc(),
+ GPU::normalizeHIPTarget(target),
+ targetFeatures))) {
return failure();
}
return success();
@@ -425,13 +442,13 @@
if (targetOptions.useAmdgcnSpirv) {
format = "rocm-spirv-fb";
} else if (deviceID == "amdgpu") {
- FailureOr<std::string> targetID =
- buildAMDGPUTargetID(b.getUnknownLoc(), targetOptions.target,
- targetOptions.targetFeatures);
+ auto targetID = parseAMDGPUTargetID(
+ b.getUnknownLoc(), GPU::normalizeHIPTarget(targetOptions.target),
+ targetOptions.targetFeatures);
if (failed(targetID)) {
return nullptr;
}
- format = *targetID;
+ format = targetID->getCanonicalFeatureString();
} else {
format = "rocm-hsaco-fb"; // legacy HIP
}
@@ -708,6 +725,12 @@
preferredSubgroupSize = attr.getPreferredSubgroupSize();
}
+ auto targetID =
+ parseAMDGPUTargetID(variantOp.getLoc(), targetArch, targetFeatures);
+ if (failed(targetID)) {
+ return failure();
+ }
+
// We name our files after the executable name so that they are easy to
// track both during compilation (logs/artifacts/etc), as outputs (final
// intermediate code/binary files), and at runtime (loaded
@@ -805,7 +828,7 @@
std::unique_ptr<llvm::TargetMachine> targetMachine;
bool isWave64 = true;
{
- llvm::Triple triple("amdgcn-amd-amdhsa");
+ llvm::Triple triple = getAMDGPUTargetTriple(targetArch);
std::string error;
const llvm::Target *target =
llvm::TargetRegistry::lookupTarget("", triple, error);
@@ -850,10 +873,6 @@
features.emplace_back("-fma-mix-insts");
}
- if (!targetFeatures.empty()) {
- features.emplace_back(targetFeatures.str());
- }
-
std::string featureStr = llvm::join(features, ",");
targetMachine.reset(target->createTargetMachine(
@@ -884,6 +903,21 @@
llvmModule->addModuleFlag(llvm::Module::Error,
"amdhsa_code_object_version", abiVersion);
+ // XNACK and SRAM ECC describe the whole code object. LLVM requires
+ // module flags for explicit settings, not subtarget features. Leave
+ // unspecified/unsupported settings absent to preserve their defaults.
+ auto addTargetIDFlag = [&](StringRef name,
+ llvm::AMDGPU::TargetIDSetting setting) {
+ using llvm::AMDGPU::TargetIDSetting;
+ if (setting == TargetIDSetting::On ||
+ setting == TargetIDSetting::Off) {
+ llvmModule->addModuleFlag(llvm::Module::Error, name,
+ setting == TargetIDSetting::On);
+ }
+ };
+ addTargetIDFlag("amdgpu.xnack", targetID->getXnackSetting());
+ addTargetIDFlag("amdgpu.sramecc", targetID->getSramEccSetting());
+
// Set the buffer OOB mode to "relaxed", since IREE does its own, more
// precise, mitigations for partially OOB buffer reads of underalligned
// vectors.
@@ -953,7 +987,7 @@
// For example 'gfx942'.
StringRef targetCPU = targetMachine->getTargetCPU();
- // For example 'amdgcn-amd-amdhsa'.
+ // For example 'amdgpu9.42-amd-amdhsa'.
std::string targetTriple = targetMachine->getTargetTriple().str();
// Run LLVM optimization passes.
@@ -1098,6 +1132,19 @@
return failure();
}
+ // LLVM a86b58566a25 started encoding gfx1250's hardwired XNACK as
+ // XNACK_ON in ELF e_flags. ROCm then looks for a gfx1250:xnack+
+ // target, which the runtime does not register. Keep the pre-existing
+ // encoding for this non-selectable feature until LLVM's ELF emission
+ // is corrected. This does not change the kernel's XNACK codegen.
+ if (targetArch == "gfx1250") {
+ char *flags =
+ targetHSACO.data() + offsetof(llvm::ELF::Elf64_Ehdr, e_flags);
+ llvm::support::endian::write32le(
+ flags, llvm::support::endian::read32le(flags) &
+ ~llvm::ELF::EF_AMDGPU_FEATURE_XNACK_V4);
+ }
+
if (targetOptions.enableRegSpillWarning) {
checkRegisterSpilling(variantOp, targetObj);
}
@@ -1131,14 +1178,12 @@
break;
}
case ContainerType::AMDGPU: {
- FailureOr<std::string> targetID =
- buildAMDGPUTargetID(variantOp.getLoc(), targetArch, targetFeatures);
- if (failed(targetID)) {
- return failure();
- }
- executableBinaryFormat = executableBuilder.getStringAttr(*targetID);
+ std::string canonicalTargetID = targetID->getCanonicalFeatureString();
+ executableBinaryFormat =
+ executableBuilder.getStringAttr(canonicalTargetID);
binaryContainer = serializeAMDGPUBinaryContainer(
- serializationOptions, variantOp, exportOps, *targetID, targetHSACO);
+ serializationOptions, variantOp, exportOps, canonicalTargetID,
+ targetHSACO);
break;
}
case ContainerType::HIP: {
diff --git a/compiler/plugins/target/ROCM/test/CMakeLists.txt b/compiler/plugins/target/ROCM/test/CMakeLists.txt
index ba3403f..eb37910 100644
--- a/compiler/plugins/target/ROCM/test/CMakeLists.txt
+++ b/compiler/plugins/target/ROCM/test/CMakeLists.txt
@@ -52,6 +52,7 @@
"smoketest.mlir"
"smoketest_hsaco.mlir"
"target_device_features.mlir"
+ "target_feature_serialization.mlir"
TOOLS
FileCheck
iree-opt
diff --git a/compiler/plugins/target/ROCM/test/default_subgroup_size.mlir b/compiler/plugins/target/ROCM/test/default_subgroup_size.mlir
index 7c8c553..b5e4309 100644
--- a/compiler/plugins/target/ROCM/test/default_subgroup_size.mlir
+++ b/compiler/plugins/target/ROCM/test/default_subgroup_size.mlir
@@ -1,15 +1,39 @@
// RUN: rm -rf %t && mkdir -p %t
// RUN: iree-opt --iree-hal-transformation-pipeline --iree-hal-target-device=hip \
-// RUN: --iree-rocm-target=gfx1250 --iree-hal-dump-executable-intermediates-to=%t %s -o /dev/null
+// RUN: --iree-rocm-target=gfx1250 --iree-hal-dump-executable-intermediates-to=%t \
+// RUN: --iree-rocm-container-type=hsaco %s | FileCheck %s --check-prefix=ELF
+// RUN: cat %t/*.rocmasm | FileCheck %s
+// RUN: rm -rf %t && mkdir -p %t
+// RUN: iree-opt --iree-hal-transformation-pipeline --iree-hal-target-device=hip \
+// RUN: --iree-rocm-target=cdna5 --iree-hal-dump-executable-intermediates-to=%t \
+// RUN: --iree-rocm-container-type=hsaco %s | FileCheck %s --check-prefix=ELF
+// RUN: cat %t/*.rocmasm | FileCheck %s
+// RUN: rm -rf %t && mkdir -p %t
+// RUN: iree-opt --iree-hal-transformation-pipeline --iree-hal-target-device=hip \
+// RUN: --iree-rocm-target=mi455x --iree-hal-dump-executable-intermediates-to=%t \
+// RUN: --iree-rocm-container-type=hsaco %s | FileCheck %s --check-prefix=ELF
// RUN: cat %t/*.rocmasm | FileCheck %s
// Verify that serialization uses the target's preferred subgroup size when an
-// export does not specify one. gfx1250 only supports wave32; asking LLVM for
+// export does not specify one. CDNA5 only supports wave32; asking LLVM for
// wave64 silently produces an executable with no kernels.
+// The ELF target ID must not require selectable XNACK on CDNA5. Check the
+// little-endian e_flags at byte offset 48: gfx1250 (0x49), SRAMECC any
+// (0x400), and XNACK unsupported (0x0).
+// ELF: data = dense<"0x7F454C46{{([0-9A-F]{88})}}49040000
+
+// Targets with selectable XNACK must retain their feature bits: gfx942
+// (0x4c), SRAMECC any (0x400), and XNACK any (0x100).
+// RUN: iree-opt --iree-hal-transformation-pipeline --iree-hal-target-device=hip \
+// RUN: --iree-rocm-target=gfx942 --iree-rocm-container-type=hsaco %s | FileCheck %s --check-prefix=ELF-CDNA3
+// ELF-CDNA3: data = dense<"0x7F454C46{{([0-9A-F]{88})}}4C050000
+
+// CHECK: .amdgcn_target "amdgpu12.50-amd-amdhsa-unknown-gfx1250"
// CHECK: .globl empty
// CHECK: amdhsa.kernels:
// CHECK: .name: empty
+// CHECK: .wavefront_size: 32
#pipeline_layout = #hal.pipeline.layout<bindings = []>
hal.executable.source public @exe {
diff --git a/compiler/plugins/target/ROCM/test/module_target_triple.mlir b/compiler/plugins/target/ROCM/test/module_target_triple.mlir
index 5ef20c2..e998028 100644
--- a/compiler/plugins/target/ROCM/test/module_target_triple.mlir
+++ b/compiler/plugins/target/ROCM/test/module_target_triple.mlir
@@ -12,7 +12,7 @@
// Therefore, the ROCMTargetBackend validation checks for the target triple to
// be set and emits an explicit error if that is not the case.
-// CHECK: target triple = "amdgcn-amd-amdhsa"
+// CHECK: target triple = "amdgpu9.42-amd-amdhsa"
#pipeline_layout = #hal.pipeline.layout<bindings = []>
hal.executable.source public @exe {
diff --git a/compiler/plugins/target/ROCM/test/target_device_features.mlir b/compiler/plugins/target/ROCM/test/target_device_features.mlir
index 7c21c47..c8ce50b 100644
--- a/compiler/plugins/target/ROCM/test/target_device_features.mlir
+++ b/compiler/plugins/target/ROCM/test/target_device_features.mlir
@@ -64,7 +64,25 @@
//
// RUN: iree-opt --pass-pipeline='builtin.module(iree-hal-assign-target-devices{targetDevices=hip},iree-hal-transformation-pipeline{serialize-executables=false})' \
// RUN: --iree-rocm-target=gfx1250 %s | FileCheck %s --check-prefixes=GFX1250
+// RUN: iree-opt --pass-pipeline='builtin.module(iree-hal-assign-target-devices{targetDevices=hip},iree-hal-transformation-pipeline{serialize-executables=false})' \
+// RUN: --iree-rocm-target=cdna5 %s | FileCheck %s --check-prefixes=GFX1250
+// RUN: iree-opt --pass-pipeline='builtin.module(iree-hal-assign-target-devices{targetDevices=hip},iree-hal-transformation-pipeline{serialize-executables=false})' \
+// RUN: --iree-rocm-target=mi455x %s | FileCheck %s --check-prefixes=GFX1250,MI455X
+// RUN: iree-opt --pass-pipeline='builtin.module(iree-hal-assign-target-devices{targetDevices=hip},iree-hal-transformation-pipeline{serialize-executables=false})' \
+// RUN: --iree-rocm-target=MI455X %s | FileCheck %s --check-prefixes=GFX1250,MI455X
+//
+// RUN: iree-opt --pass-pipeline='builtin.module(iree-hal-assign-target-devices{targetDevices=amdgpu},iree-hal-transformation-pipeline{serialize-executables=false})' \
+// RUN: --iree-rocm-target=gfx1250 %s | FileCheck %s --check-prefixes=GFX1250,AMDGPU
+// RUN: iree-opt --pass-pipeline='builtin.module(iree-hal-assign-target-devices{targetDevices=amdgpu},iree-hal-transformation-pipeline{serialize-executables=false})' \
+// RUN: --iree-rocm-target=cdna5 %s | FileCheck %s --check-prefixes=GFX1250,AMDGPU
+// RUN: iree-opt --pass-pipeline='builtin.module(iree-hal-assign-target-devices{targetDevices=amdgpu},iree-hal-transformation-pipeline{serialize-executables=false})' \
+// RUN: --iree-rocm-target=mi455x %s | FileCheck %s --check-prefixes=GFX1250,AMDGPU
+// RUN: iree-opt --pass-pipeline='builtin.module(iree-hal-assign-target-devices{targetDevices=amdgpu},iree-hal-transformation-pipeline{serialize-executables=false})' \
+// RUN: --iree-rocm-target=MI455X %s | FileCheck %s --check-prefixes=GFX1250,AMDGPU
+// RUN: iree-opt --pass-pipeline='builtin.module(iree-hal-assign-target-devices{targetDevices=amdgpu},iree-hal-transformation-pipeline{serialize-executables=false})' \
+// RUN: --iree-rocm-target=mi300x --iree-rocm-target-features=+sramecc,-xnack %s | FileCheck %s --check-prefixes=GFX942,AMDGPU-FEATURES
+// AMDGPU-FEATURES: #hal.executable.target<"rocm", "gfx942:sramecc+:xnack-",
// GFX942: target_info = #iree_gpu.target<arch = "gfx942",
// GFX942-SAME: wgp = <compute = fp64|fp32|fp16|int64|int32|int16|int8, storage = b64|b32|b16|b8,
// GFX942-SAME: subgroup = shuffle|arithmetic, dot = dp4xi8toi32,
@@ -128,8 +146,11 @@
// RX9070: chip = <wgp_count = 28, sku = "rx9070", memory_bandwidth_tbps = 6.400000e-01 : f32, perf_tflops = {fp16 = 1.450000e+02 : f32, fp32 = 3.610000e+01 : f32, fp8 = 2.890000e+02 : f32, int8 = 2.890000e+02 : f32}>>
// R9700: chip = <wgp_count = 32, sku = "r9700", memory_bandwidth_tbps = 6.400000e-01 : f32, perf_tflops = {fp16 = 1.910000e+02 : f32, fp32 = 4.780000e+01 : f32, fp8 = 3.830000e+02 : f32, int8 = 3.830000e+02 : f32}>>
-// Note: The gfx1250 target is experimental and contains placeholder values.
+// CDNA5 architecture limits and MI455X product specifications.
+// AMDGPU: #hal.executable.target<"rocm", "gfx1250",
// GFX1250: target_info = #iree_gpu.target<arch = "gfx1250",
+// GFX1250-SAME: wgp = <compute = fp64|fp32|fp16|int64|int32|int16|int8, storage = b64|b32|b16|b8,
+// GFX1250-SAME: subgroup = shuffle|arithmetic,
// GFX1250-SAME: mma = [<WMMA_F32_16x16x4_F32>,
// GFX1250-SAME: <WMMA_F32_16x16x32_F16>, <WMMA_F32_16x16x32_BF16>, <WMMA_F16_16x16x32_F16>, <WMMA_BF16_16x16x32_BF16>,
// GFX1250-SAME: <WMMA_F32_16x16x64_F8E4M3FN>, <WMMA_F32_16x16x64_F8E4M3FN_F8E5M2>, <WMMA_F32_16x16x64_F8E5M2>, <WMMA_F32_16x16x64_F8E5M2_F8E4M3FN>,
@@ -137,8 +158,11 @@
// GFX1250-SAME: <WMMA_I32_16x16x64_I8>,
// GFX1250-SAME: <WMMA_F32_16x16x128_F8E5M2>, <WMMA_F32_16x16x128_F8E5M2_F8E4M3FN>, <WMMA_F32_16x16x128_F8E4M3FN>, <WMMA_F32_16x16x128_F8E4M3FN_F8E5M2>,
// GFX1250-SAME: <WMMA_F16_16x16x128_F8E5M2>, <WMMA_F16_16x16x128_F8E5M2_F8E4M3FN>, <WMMA_F16_16x16x128_F8E4M3FN>, <WMMA_F16_16x16x128_F8E4M3FN_F8E5M2>]
-// GFX1250-SAME: subgroup_size_choices = [32]
+// GFX1250-SAME: subgroup_size_choices = [32], max_workgroup_sizes = [1024, 1024, 1024],
+// GFX1250-SAME: max_thread_count_per_workgroup = 1024, max_workgroup_memory_bytes = 327680,
+// GFX1250-SAME: max_workgroup_counts = [2147483647, 2147483647, 2147483647],
// GFX1250-SAME: max_load_instruction_bits = 128, simds_per_wgp = 4, vgpr_space_bits = 32768, workgroup_memory_bank_count = 64
+// MI455X: chip = <wgp_count = 256, sku = "mi455x", memory_bandwidth_tbps = 2.330000e+01 : f32, perf_tflops = {fp16 = 5.030000e+03 : f32, fp32 = 3.150000e+02 : f32, fp4 = 4.026000e+04 : f32, fp6 = 2.013000e+04 : f32, fp8 = 2.013000e+04 : f32}>>
stream.executable public @reduce_dispatch {
stream.executable.export @reduce_dispatch workgroups() -> (index, index, index) {
diff --git a/compiler/plugins/target/ROCM/test/target_feature_serialization.mlir b/compiler/plugins/target/ROCM/test/target_feature_serialization.mlir
new file mode 100644
index 0000000..f3b452c
--- /dev/null
+++ b/compiler/plugins/target/ROCM/test/target_feature_serialization.mlir
@@ -0,0 +1,80 @@
+// XNACK and SRAM ECC must reach LLVM as module flags and survive native
+// serialization in the ELF target ID. Omitting a mode preserves Any.
+// The LLVM target machine must not receive these modes as subtarget features.
+// ASM: ; llc -mtriple=amdgpu9.42-amd-amdhsa -mcpu=gfx942 -mattr='-fma-mix-insts'
+
+// RUN: rm -rf %t && mkdir -p %t
+// RUN: iree-opt --iree-hal-transformation-pipeline --iree-hal-target-device=hip \
+// RUN: --iree-rocm-target=mi300x \
+// RUN: --iree-hal-dump-executable-intermediates-to=%t --iree-rocm-container-type=hsaco %s | FileCheck %s --check-prefix=ELF-ANY
+// RUN: cat %t/*.linked.ll | FileCheck %s --check-prefix=IR-ANY --implicit-check-not='!"amdgpu.xnack"' --implicit-check-not='!"amdgpu.sramecc"'
+// RUN: cat %t/*.rocmasm | FileCheck %s --check-prefix=ASM
+// ELF-ANY: data = dense<"0x7F454C46{{([0-9A-F]{88})}}4C050000
+// IR-ANY: target triple = "amdgpu9.42-amd-amdhsa"
+
+// RUN: rm -rf %t && mkdir -p %t
+// RUN: iree-opt --iree-hal-transformation-pipeline --iree-hal-target-device=hip \
+// RUN: --iree-rocm-target=mi300x --iree-rocm-target-features=+xnack \
+// RUN: --iree-hal-dump-executable-intermediates-to=%t --iree-rocm-container-type=hsaco %s | FileCheck %s --check-prefix=ELF-XNACKON
+// RUN: cat %t/*.linked.ll | FileCheck %s --check-prefix=IR-XNACKON --implicit-check-not='!"amdgpu.sramecc"'
+// RUN: cat %t/*.rocmasm | FileCheck %s --check-prefix=ASM
+// ELF-XNACKON: data = dense<"0x7F454C46{{([0-9A-F]{88})}}4C070000
+// IR-XNACKON: target triple = "amdgpu9.42-amd-amdhsa"
+// IR-XNACKON-DAG: !{i32 1, !"amdgpu.xnack", i32 1}
+
+// RUN: rm -rf %t && mkdir -p %t
+// RUN: iree-opt --iree-hal-transformation-pipeline --iree-hal-target-device=hip \
+// RUN: --iree-rocm-target=mi300x --iree-rocm-target-features=-xnack \
+// RUN: --iree-hal-dump-executable-intermediates-to=%t --iree-rocm-container-type=hsaco %s | FileCheck %s --check-prefix=ELF-XNACKOFF
+// RUN: cat %t/*.linked.ll | FileCheck %s --check-prefix=IR-XNACKOFF --implicit-check-not='!"amdgpu.sramecc"'
+// RUN: cat %t/*.rocmasm | FileCheck %s --check-prefix=ASM
+// ELF-XNACKOFF: data = dense<"0x7F454C46{{([0-9A-F]{88})}}4C060000
+// IR-XNACKOFF: target triple = "amdgpu9.42-amd-amdhsa"
+// IR-XNACKOFF-DAG: !{i32 1, !"amdgpu.xnack", i32 0}
+
+// RUN: rm -rf %t && mkdir -p %t
+// RUN: iree-opt --iree-hal-transformation-pipeline --iree-hal-target-device=hip \
+// RUN: --iree-rocm-target=mi300x --iree-rocm-target-features=+sramecc,-xnack \
+// RUN: --iree-hal-dump-executable-intermediates-to=%t --iree-rocm-container-type=hsaco %s | FileCheck %s --check-prefix=ELF-ECCON
+// RUN: cat %t/*.linked.ll | FileCheck %s --check-prefix=IR-ECCON
+// RUN: cat %t/*.rocmasm | FileCheck %s --check-prefix=ASM
+// ELF-ECCON: data = dense<"0x7F454C46{{([0-9A-F]{88})}}4C0E0000
+// IR-ECCON: target triple = "amdgpu9.42-amd-amdhsa"
+// IR-ECCON-DAG: !{i32 1, !"amdgpu.sramecc", i32 1}
+// IR-ECCON-DAG: !{i32 1, !"amdgpu.xnack", i32 0}
+
+// RUN: rm -rf %t && mkdir -p %t
+// RUN: iree-opt --iree-hal-transformation-pipeline --iree-hal-target-device=hip \
+// RUN: --iree-rocm-target=mi300x --iree-rocm-target-features=-sramecc,+xnack \
+// RUN: --iree-hal-dump-executable-intermediates-to=%t --iree-rocm-container-type=hsaco %s | FileCheck %s --check-prefix=ELF-ECCOFF
+// RUN: cat %t/*.linked.ll | FileCheck %s --check-prefix=IR-ECCOFF
+// RUN: cat %t/*.rocmasm | FileCheck %s --check-prefix=ASM
+// ELF-ECCOFF: data = dense<"0x7F454C46{{([0-9A-F]{88})}}4C0B0000
+// IR-ECCOFF: target triple = "amdgpu9.42-amd-amdhsa"
+// IR-ECCOFF-DAG: !{i32 1, !"amdgpu.sramecc", i32 0}
+// IR-ECCOFF-DAG: !{i32 1, !"amdgpu.xnack", i32 1}
+
+// RUN: not iree-opt --iree-hal-transformation-pipeline --iree-hal-target-device=hip \
+// RUN: --iree-rocm-target=mi455x --iree-rocm-target-features=-xnack %s -o /dev/null 2>&1 | FileCheck %s --check-prefix=CDNA5OFF
+// CDNA5OFF: invalid ROCM target ID 'gfx1250:xnack-'
+
+// RUN: not iree-opt --iree-hal-transformation-pipeline --iree-hal-target-device=hip \
+// RUN: --iree-rocm-target=cdna5 --iree-rocm-target-features=+xnack %s -o /dev/null 2>&1 | FileCheck %s --check-prefix=CDNA5ON
+// CDNA5ON: invalid ROCM target ID 'gfx1250:xnack+'
+
+// RUN: not iree-opt --iree-hal-transformation-pipeline --iree-hal-target-device=hip \
+// RUN: --iree-rocm-target=gfx1100 --iree-rocm-target-features=+sramecc %s -o /dev/null 2>&1 | FileCheck %s --check-prefix=INVALIDFEATURE
+// INVALIDFEATURE: invalid ROCM target ID 'gfx1100:sramecc+'
+
+#pipeline_layout = #hal.pipeline.layout<bindings = []>
+hal.executable.source public @exe {
+ hal.executable.export public @empty ordinal(0) layout(#pipeline_layout) count(%arg0: !hal.device) -> (index, index, index) {
+ %c1 = arith.constant 1 : index
+ hal.return %c1, %c1, %c1 : index, index, index
+ } attributes {workgroup_size = [1 : index, 1 : index, 1 : index]}
+ builtin.module {
+ llvm.func @empty() attributes {rocdl.kernel} {
+ llvm.return
+ }
+ }
+}
diff --git a/compiler/src/iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUEnums.td b/compiler/src/iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUEnums.td
index 80f6311..59348ec 100644
--- a/compiler/src/iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUEnums.td
+++ b/compiler/src/iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUEnums.td
@@ -148,7 +148,7 @@
// * 3 = CDNA4
// * 8 = RDNA3
// * 9 = RDNA4
-// * A = gfx1250
+// * A = CDNA5
// * C = element type of A-matrix:
// * 0 = 64-bit float (e.g. IEEE754 double precision)
// * 1 = 32-bit float (e.g. IEEE754 single precision, and "xf32" fast variants)
@@ -241,7 +241,7 @@
def WMMAR4_F32_16x16x16_F8E4M3FN_F8E5M2 : I32EnumAttrCase<"WMMAR4_F32_16x16x16_F8E4M3FN_F8E5M2", 0x1933>;
def WMMAR4_I32_16x16x16_I8 : I32EnumAttrCase<"WMMAR4_I32_16x16x16_I8", 0x19C0>;
-// Introduced in gfx1250.
+// Introduced in CDNA5.
// K=4.
def WMMA_F32_16x16x4_F32 : I32EnumAttrCase<"WMMA_F32_16x16x4_F32", 0x1A10>;
// K=32.
@@ -366,7 +366,7 @@
WMMAR4_F32_16x16x16_F8E4M3FN_F8E5M2,
WMMAR4_I32_16x16x16_I8,
- // gfx1250 intrinsics.
+ // CDNA5 intrinsics.
// K=4.
WMMA_F32_16x16x4_F32,
// K=32.
diff --git a/compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/KnownTargets.cpp b/compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/KnownTargets.cpp
index 65edff2..1d936da 100644
--- a/compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/KnownTargets.cpp
+++ b/compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/KnownTargets.cpp
@@ -210,6 +210,63 @@
// to 64-bits.
//===----------------------------------------------------------------------===//
+// CDNA5 WGP limits: sections 1 and 3 of the architecture white paper and
+// sections 3.2, 3.3.2, and 11.1 of the ISA reference.
+// https://www.amd.com/content/dam/amd/en/documents/products/technologies/cdna/amd-cdna5-whitepaper.pdf
+// https://www.amd.com/content/dam/amd/en/documents/instinct-tech-docs/instruction-set-architectures/amd-instinct-cdna5-instruction-set-architecture.pdf
+const WgpDetails *getCDNA5WgpDetails() {
+ static const MMAIntrinsic cdna5MMAOps[] = {
+ // K=4.
+ MMAIntrinsic::WMMA_F32_16x16x4_F32,
+ // K=32.
+ MMAIntrinsic::WMMA_F32_16x16x32_F16,
+ MMAIntrinsic::WMMA_F32_16x16x32_BF16,
+ MMAIntrinsic::WMMA_F16_16x16x32_F16,
+ MMAIntrinsic::WMMA_BF16_16x16x32_BF16,
+ // K=64.
+ MMAIntrinsic::WMMA_F32_16x16x64_F8E4M3FN,
+ MMAIntrinsic::WMMA_F32_16x16x64_F8E4M3FN_F8E5M2,
+ MMAIntrinsic::WMMA_F32_16x16x64_F8E5M2,
+ MMAIntrinsic::WMMA_F32_16x16x64_F8E5M2_F8E4M3FN,
+ MMAIntrinsic::WMMA_F16_16x16x64_F8E4M3FN,
+ MMAIntrinsic::WMMA_F16_16x16x64_F8E4M3FN_F8E5M2,
+ MMAIntrinsic::WMMA_F16_16x16x64_F8E5M2,
+ MMAIntrinsic::WMMA_F16_16x16x64_F8E5M2_F8E4M3FN,
+ MMAIntrinsic::WMMA_I32_16x16x64_I8,
+ // K=128.
+ MMAIntrinsic::WMMA_F32_16x16x128_F8E5M2,
+ MMAIntrinsic::WMMA_F32_16x16x128_F8E5M2_F8E4M3FN,
+ MMAIntrinsic::WMMA_F32_16x16x128_F8E4M3FN,
+ MMAIntrinsic::WMMA_F32_16x16x128_F8E4M3FN_F8E5M2,
+ MMAIntrinsic::WMMA_F16_16x16x128_F8E5M2,
+ MMAIntrinsic::WMMA_F16_16x16x128_F8E5M2_F8E4M3FN,
+ MMAIntrinsic::WMMA_F16_16x16x128_F8E4M3FN,
+ MMAIntrinsic::WMMA_F16_16x16x128_F8E4M3FN_F8E5M2,
+ };
+
+ // Scaled WMMA and asynchronous global-to-LDS lowering are not enabled yet.
+ static const WgpDetails cdna5Wgp = {allComputeBits,
+ allStorageBits,
+ allSubgroupOps,
+ DotProductOps::None,
+ /*mmaCount=*/std::size(cdna5MMAOps),
+ /*mmaOps=*/cdna5MMAOps,
+ /*scaledMmaCount=*/0,
+ /*scaledMmaOps=*/nullptr,
+ {32, 32},
+ {1024, 1024, 1024},
+ 1024,
+ 320 * 1024,
+ {0x7fffffff, 0x7fffffff, 0x7fffffff},
+ /*maxLoadInstructionBits=*/128,
+ /*simdsPerWgp=*/4,
+ // Up to 1024 32-bit VGPRs per thread.
+ /*vgprSpaceBits=*/1024 * 32,
+ /*dmaSizes=*/std::nullopt,
+ /*workgroupMemoryBankCount=*/64};
+ return &cdna5Wgp;
+}
+
const WgpDetails *getCDNA4WgpDetails() {
static const MMAIntrinsic cdna4MMAOps[] = {
// Introduced in CDNA4
@@ -539,60 +596,8 @@
return &rdna1Wgp;
}
-// Experimental gfx1250 WGP details.
-const WgpDetails *getGfx1250WgpDetails() {
- static const MMAIntrinsic gfx1250MMAOps[] = {
- // K=4.
- MMAIntrinsic::WMMA_F32_16x16x4_F32,
- // K=32.
- MMAIntrinsic::WMMA_F32_16x16x32_F16,
- MMAIntrinsic::WMMA_F32_16x16x32_BF16,
- MMAIntrinsic::WMMA_F16_16x16x32_F16,
- MMAIntrinsic::WMMA_BF16_16x16x32_BF16,
- // K=64.
- MMAIntrinsic::WMMA_F32_16x16x64_F8E4M3FN,
- MMAIntrinsic::WMMA_F32_16x16x64_F8E4M3FN_F8E5M2,
- MMAIntrinsic::WMMA_F32_16x16x64_F8E5M2,
- MMAIntrinsic::WMMA_F32_16x16x64_F8E5M2_F8E4M3FN,
- MMAIntrinsic::WMMA_F16_16x16x64_F8E4M3FN,
- MMAIntrinsic::WMMA_F16_16x16x64_F8E4M3FN_F8E5M2,
- MMAIntrinsic::WMMA_F16_16x16x64_F8E5M2,
- MMAIntrinsic::WMMA_F16_16x16x64_F8E5M2_F8E4M3FN,
- MMAIntrinsic::WMMA_I32_16x16x64_I8,
- // K=128.
- MMAIntrinsic::WMMA_F32_16x16x128_F8E5M2,
- MMAIntrinsic::WMMA_F32_16x16x128_F8E5M2_F8E4M3FN,
- MMAIntrinsic::WMMA_F32_16x16x128_F8E4M3FN,
- MMAIntrinsic::WMMA_F32_16x16x128_F8E4M3FN_F8E5M2,
- MMAIntrinsic::WMMA_F16_16x16x128_F8E5M2,
- MMAIntrinsic::WMMA_F16_16x16x128_F8E5M2_F8E4M3FN,
- MMAIntrinsic::WMMA_F16_16x16x128_F8E4M3FN,
- MMAIntrinsic::WMMA_F16_16x16x128_F8E4M3FN_F8E5M2,
- };
-
- static const WgpDetails gfx1250Wgp = {allComputeBits,
- allStorageBits,
- allSubgroupOps,
- DotProductOps::None,
- /*mmaCount=*/std::size(gfx1250MMAOps),
- /*mmaOps=*/gfx1250MMAOps,
- /*scaledMmaCount=*/0,
- /*scaledMmaOps=*/nullptr,
- {32, 32},
- {1024, 1024, 1024},
- 1024,
- 320 * 1024,
- {0x7fffffff, 0x7fffffff, 0x7fffffff},
- /*maxLoadInstructionBits=*/128,
- /*simdsPerWgp=*/4,
- // 4 banks of 256 32-bit registers.
- /*vgprSpaceBits=*/256 * 4 * 32,
- /*dmaSizes=*/std::nullopt,
- /*workgroupMemoryBankCount=*/64};
- return &gfx1250Wgp;
-}
-
std::optional<TargetDetails> getAMDGPUTargetDetails(StringRef target) {
+ const WgpDetails *cdna5Wgp = getCDNA5WgpDetails();
const WgpDetails *cdna4Wgp = getCDNA4WgpDetails();
const WgpDetails *cdna3Wgp = getCDNA3WgpDetails();
const WgpDetails *cdna2Wgp = getCDNA2WgpDetails();
@@ -601,9 +606,21 @@
const WgpDetails *rdna3Wgp = getRDNA3WgpDetails();
const WgpDetails *rdna2Wgp = getRDNA2WgpDetails();
const WgpDetails *rdna1Wgp = getRDNA1WgpDetails();
- const WgpDetails *gfx1250Wgp = getGfx1250WgpDetails(); // Experimental.
// --- CDNA --- //
+ // "AMD Instinct MI455X GPU" feature summary in the CDNA5 white paper:
+ // https://www.amd.com/content/dam/amd/en/documents/products/technologies/cdna/amd-cdna5-whitepaper.pdf
+ // Matrix rates are converted from PFLOPs/s to TFLOPs/s. The white paper
+ // does not specify an INT8 rate.
+ static const ChipDetails mi455xChip = {256,
+ "mi455x",
+ 23.3f,
+ {{ComputeBitwidths::FP32, 315.0f},
+ {ComputeBitwidths::FP16, 5030.0f},
+ {ComputeBitwidths::FP8, 20130.0f},
+ {ComputeBitwidths::FP6, 20130.0f},
+ {ComputeBitwidths::FP4, 40260.0f}}};
+
// "AMD Instinct MI350 Series Product Offerings" in Page 18 of
// https://www.amd.com/content/dam/amd/en/documents/instinct-tech-docs/white-papers/amd-cdna-4-architecture-whitepaper.pdf
static const ChipDetails mi350xChip = {256,
@@ -798,6 +815,8 @@
// See https://llvm.org/docs/AMDGPUUsage.html#processors for gfxN to
// cdnaN/rdnaN mapping.
return llvm::StringSwitch<std::optional<TargetDetails>>(target.lower())
+ .Case("mi455x", TargetDetails{cdna5Wgp, &mi455xChip})
+ .Cases({"cdna5", "gfx1250"}, TargetDetails{cdna5Wgp, nullptr})
.Case("mi355x", TargetDetails{cdna4Wgp, &mi355xChip})
.Case("mi350x", TargetDetails{cdna4Wgp, &mi350xChip})
.Cases({"cdna4", "gfx950"}, TargetDetails{cdna4Wgp, nullptr})
@@ -851,12 +870,12 @@
TargetDetails{rdna2Wgp, nullptr})
.Cases({"rdna1", "gfx1010", "gfx1011", "gfx1012", "gfx1013"},
TargetDetails{rdna1Wgp, nullptr})
- .Case("gfx1250", TargetDetails{gfx1250Wgp, nullptr})
.Default(std::nullopt);
}
StringRef normalizeAMDGPUTarget(StringRef target) {
return llvm::StringSwitch<StringRef>(target.lower())
+ .Cases({"mi455x", "cdna5", "gfx1250"}, /*Value=*/"gfx1250")
.Cases({"mi350x", "mi355x", "gfx950"}, /*Value=*/"gfx950")
.Cases({"mi300a", "mi300x", "mi308x", "mi325x", "gfx942"},
/*Value=*/"gfx942")
@@ -872,7 +891,6 @@
.Cases({"phoenix", "gfx1103"}, /*Value=*/"gfx1103")
.Cases({"strix-point", "gfx1150"}, /*Value=*/"gfx1150")
.Cases({"strix-halo", "gfx1151"}, /*Value=*/"gfx1151")
- .Case("gfx1250", /*Value=*/"gfx1250")
.Default("");
}
@@ -1645,8 +1663,8 @@
return kCDNA4Seeds;
}
- // RDNA4 is gfx1200/gfx1201 (major=12, minor<=1). Note: gfx1250 (minor=50)
- // is a separate experimental target and should not use RDNA4 seeds.
+ // RDNA4 is gfx1200/gfx1201 (major=12, minor=0). CDNA5 is gfx1250
+ // (major=12, minor=5) and should not use RDNA4 seeds.
bool isRDNA4 = succeeded(chipset) && chipset->majorVersion == 12 &&
chipset->minorVersion <= 1;
if (isRDNA4 || arch == "rdna4") {
diff --git a/docs/website/docs/guides/deployment-configurations/gpu-rocm.md b/docs/website/docs/guides/deployment-configurations/gpu-rocm.md
index 2b51e47..e3ae06f 100644
--- a/docs/website/docs/guides/deployment-configurations/gpu-rocm.md
+++ b/docs/website/docs/guides/deployment-configurations/gpu-rocm.md
@@ -195,6 +195,7 @@
| AMD MI325X | `mi325x` | `gfx942` | `cdna3` |
| AMD MI350X | `mi350x` | `gfx950` | `cdna4` |
| AMD MI355X | `mi355x` | `gfx950` | `cdna4` |
+| AMD MI455X | `mi455x` | `gfx1250` | `cdna5` |
| AMD PRO V710 | `v710` | `gfx1101` | `rdna3` |
| AMD PRO W7700 | `w7700` | `gfx1101` | `rdna3` |
| AMD PRO W7800 | `w7800` | `gfx1100` | `rdna3` |
diff --git a/runtime/src/iree/hal/drivers/amdgpu/device/binaries/README.md b/runtime/src/iree/hal/drivers/amdgpu/device/binaries/README.md
index 4efaafa..270d204 100644
--- a/runtime/src/iree/hal/drivers/amdgpu/device/binaries/README.md
+++ b/runtime/src/iree/hal/drivers/amdgpu/device/binaries/README.md
@@ -143,7 +143,7 @@
| `gfx10.3` RDNA | `gfx1030`, `gfx1031`, `gfx1032`, `gfx1033`, `gfx1034`, `gfx1035`, `gfx1036` | `gfx10-3-generic` |
| `gfx11` RDNA/APU | `gfx1100`, `gfx1101`, `gfx1102`, `gfx1103`, `gfx1150`, `gfx1151`, `gfx1152`, `gfx1153`, `gfx1170`, `gfx1171`, `gfx1172` | `gfx11-generic` |
| `gfx12` RDNA | `gfx1200`, `gfx1201` | `gfx12-generic` |
-| `gfx12.5` RDNA | `gfx1250`, `gfx1251` | `gfx12-5-generic` |
+| `gfx12.5` CDNA5 | `gfx1250`, `gfx1251` | `gfx12-5-generic` |
`gfx12-5-generic` is available as an explicit selector, but it is intentionally
not part of the default checked-in prebuilt set until ROCm/LLVM can link these
