[SPIRV][GPU] Add Mali G720 and other Arm 5 Gen target details (#24595)
Based on the official `vulkaninfo`, define `arm5gen` as an IREE Vulkan
target for devices such as Mali G720 Immortalis. Include all compute
bitwidths into the declaration, based on the official support for
`shaderInt64` capability - besides, initial experiments on a set of
Torch and TFLite models have shown that Arm driver emulation for wide
integers holds up against the current emulation in IREE.
Signed-off-by: Artem Gindinson <gindinson@roofline.ai>
diff --git a/compiler/plugins/target/VulkanSPIRV/test/target_device_features.mlir b/compiler/plugins/target/VulkanSPIRV/test/target_device_features.mlir
index 8a35706..1b20320 100644
--- a/compiler/plugins/target/VulkanSPIRV/test/target_device_features.mlir
+++ b/compiler/plugins/target/VulkanSPIRV/test/target_device_features.mlir
@@ -28,6 +28,15 @@
// RUN: iree-opt --pass-pipeline='builtin.module(iree-hal-assign-target-devices{targetDevices=vulkan},iree-hal-transformation-pipeline{serialize-executables=false})' \
// RUN: --iree-vulkan-target=mali-g615 %s | FileCheck %s --check-prefixes=VALHALL,VALHALL4
+// RUN: iree-opt --pass-pipeline='builtin.module(iree-hal-assign-target-devices{targetDevices=vulkan},iree-hal-transformation-pipeline{serialize-executables=false})' \
+// RUN: --iree-vulkan-target=arm5gen %s | FileCheck %s --check-prefixes=ARM5GEN
+// RUN: iree-opt --pass-pipeline='builtin.module(iree-hal-assign-target-devices{targetDevices=vulkan},iree-hal-transformation-pipeline{serialize-executables=false})' \
+// RUN: --iree-vulkan-target=mali-g720 %s | FileCheck %s --check-prefixes=ARM5GEN
+// RUN: iree-opt --pass-pipeline='builtin.module(iree-hal-assign-target-devices{targetDevices=vulkan},iree-hal-transformation-pipeline{serialize-executables=false})' \
+// RUN: --iree-vulkan-target=mali-g725 %s | FileCheck %s --check-prefixes=ARM5GEN
+// RUN: iree-opt --pass-pipeline='builtin.module(iree-hal-assign-target-devices{targetDevices=vulkan},iree-hal-transformation-pipeline{serialize-executables=false})' \
+// RUN: --iree-vulkan-target=mali-g1 %s | FileCheck %s --check-prefixes=ARM5GEN
+
// VALHALL1: target_info = #iree_gpu.target<arch = "valhall1",
// VALHALL2: target_info = #iree_gpu.target<arch = "valhall2",
// VALHALL3: target_info = #iree_gpu.target<arch = "valhall3",
@@ -47,6 +56,22 @@
// VALHALL-SAME: max_workgroup_memory_bytes = 32768
// VALHALL-SAME: max_workgroup_counts = [65535, 65535, 65535]
+// ARM5GEN: target_info = #iree_gpu.target<arch = "arm5gen",
+// ARM5GEN-SAME: features = "spirv:v1.6,cap:Shader"
+// ARM5GEN-SAME: wgp = <compute =
+// ARM5GEN-SAME: fp64|fp32|fp16|int64|int32|int16|int8
+// ARM5GEN-SAME: storage =
+// ARM5GEN-SAME: b64|b32|b16|b8
+// ARM5GEN-SAME: subgroup =
+// ARM5GEN-SAME: shuffle|arithmetic
+// ARM5GEN-SAME: dot =
+// ARM5GEN-SAME: dp4xi8toi32
+// ARM5GEN-SAME: subgroup_size_choices = [16]
+// ARM5GEN-SAME: max_workgroup_sizes = [1024, 1024, 1024]
+// ARM5GEN-SAME: max_thread_count_per_workgroup = 1024
+// ARM5GEN-SAME: max_workgroup_memory_bytes = 32768
+// ARM5GEN-SAME: max_workgroup_counts = [65535, 65535, 65535]
+
stream.executable public @reduce_dispatch {
stream.executable.export @reduce_dispatch workgroups() -> (index, index, index) {
%x, %y, %z = iree_tensor_ext.dispatch.workgroup_count_from_slice()
diff --git a/compiler/src/iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUAttrs.td b/compiler/src/iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUAttrs.td
index 75e826f..724d541 100644
--- a/compiler/src/iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUAttrs.td
+++ b/compiler/src/iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUAttrs.td
@@ -918,7 +918,7 @@
return getArch().starts_with("gfx") || getArch().starts_with("rdna");
}
bool isApple() const { return getArch().starts_with("apple"); }
- bool isARM() const { return getArch().starts_with("valhall"); }
+ bool isARM() const { return getArch().starts_with("valhall") || getArch().starts_with("arm"); }
bool isBroadcom() const { return getArch().starts_with("videocore_"); }
bool isNVIDIA() const { return getArch().starts_with("sm_"); }
bool isQualcomm() const { return getArch().starts_with("adreno"); }
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 a3b3bfd..680d248 100644
--- a/compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/KnownTargets.cpp
+++ b/compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/KnownTargets.cpp
@@ -900,6 +900,21 @@
// Known ARM target details
//===----------------------------------------------------------------------===//
+const WgpDetails *get5GenArmWgpDetails() {
+ ComputeBitwidths computeBitwdiths = allComputeBits;
+ // clang-format off
+ static const WgpDetails fifthGenWgp = {
+ computeBitwdiths, allStorageBits, allSubgroupOps, allDotProductOps,
+ /*mmaCount=*/0, /*mmaOps=*/nullptr,
+ /*scaledMmaCount=*/0, /*scaledMmaOps=*/nullptr,
+ /*subgroupSizeChoices=*/{16, 16}, /*maxWorkgroupSizes=*/{1024, 1024, 1024},
+ /*maxThreadSize=*/1024,
+ /*maxWorkgroupMemoryBytes=*/32 * 1024, // Vulkan: maxComputeSharedMemorySize
+ /*maxWorkgroupCounts=*/{0xffff, 0xffff, 0xffff}};
+ // clang-format on
+ return &fifthGenWgp;
+}
+
const WgpDetails *getValhallWgpDetails() {
// Recent drivers report support for shaderInt64. Aside from not widely
// applicable, we don't know whether that's emulated and how performant it is.
@@ -927,12 +942,16 @@
std::optional<TargetDetails> getARMGPUTargetDetails(StringRef target) {
const WgpDetails *valhallWgp = getValhallWgpDetails();
+ const WgpDetails *fifthGenArmWgp = get5GenArmWgpDetails();
// Note that the underlying GPU may have certain capabilities but the Android
// version and driver stack may not expose them. So the following is just and
// will always be approximate.
return llvm::StringSwitch<std::optional<TargetDetails>>(target.lower())
+ // Mali-G720: https://vulkan.gpuinfo.org/displayreport.php?id=36655
+ .Cases({"mali-g720", "mali-g725", "mali-g1", "arm5gen"},
+ TargetDetails{fifthGenArmWgp, nullptr})
// Mali-G715: https://vulkan.gpuinfo.org/displayreport.php?id=29754
.Cases({"mali-g715", "mali-g615", "valhall4"},
TargetDetails{valhallWgp, nullptr})
@@ -956,6 +975,7 @@
}
return llvm::StringSwitch<StringRef>(target.lower())
+ .Cases({"mali-g720", "mali-g725", "mali-g1", "arm5gen"}, "arm5gen")
.Cases({"mali-g715", "mali-g615"}, "valhall4")
.Cases({"mali-g710", "mali-g510", "mali-g310"}, "valhall3")
.Case("mali-g78", "valhall2")
