[Vulkan][SPIRV] Add initial LIT for known Vulkan target features (Arm GPUs)
Create a regression test for Arm Mali targets, similar to CUDA/ROCM
tests for recorded target capabilities. This can be later expanded for
all known Vulkan targets within IREE.
Additionally, fix the lexeme for Mali G78.
Assisted-by: Codex <noreply@openai.com>
Signed-off-by: Artem Gindinson <gindinson@roofline.ai>
diff --git a/compiler/plugins/target/VulkanSPIRV/test/BUILD.bazel b/compiler/plugins/target/VulkanSPIRV/test/BUILD.bazel
index cbc9a61..7afbb36 100644
--- a/compiler/plugins/target/VulkanSPIRV/test/BUILD.bazel
+++ b/compiler/plugins/target/VulkanSPIRV/test/BUILD.bazel
@@ -20,6 +20,7 @@
"hal_target_device_attributes.mlir",
"materialize_homogeneous_encodings.mlir",
"smoketest.mlir",
+ "target_device_features.mlir",
],
include = ["*.mlir"],
),
diff --git a/compiler/plugins/target/VulkanSPIRV/test/CMakeLists.txt b/compiler/plugins/target/VulkanSPIRV/test/CMakeLists.txt
index 2a697c0..edcedf1 100644
--- a/compiler/plugins/target/VulkanSPIRV/test/CMakeLists.txt
+++ b/compiler/plugins/target/VulkanSPIRV/test/CMakeLists.txt
@@ -17,6 +17,7 @@
"hal_target_device_attributes.mlir"
"materialize_homogeneous_encodings.mlir"
"smoketest.mlir"
+ "target_device_features.mlir"
TOOLS
FileCheck
iree-compile
diff --git a/compiler/plugins/target/VulkanSPIRV/test/target_device_features.mlir b/compiler/plugins/target/VulkanSPIRV/test/target_device_features.mlir
new file mode 100644
index 0000000..8a35706
--- /dev/null
+++ b/compiler/plugins/target/VulkanSPIRV/test/target_device_features.mlir
@@ -0,0 +1,60 @@
+// 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=valhall %s | FileCheck %s --check-prefixes=VALHALL,VALHALL1
+// 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=valhall1 %s | FileCheck %s --check-prefixes=VALHALL,VALHALL1
+// 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-g77 %s | FileCheck %s --check-prefixes=VALHALL,VALHALL1
+// 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-g57 %s | FileCheck %s --check-prefixes=VALHALL,VALHALL1
+//
+// 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=valhall2 %s | FileCheck %s --check-prefixes=VALHALL,VALHALL2
+// 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-g78 %s | FileCheck %s --check-prefixes=VALHALL,VALHALL2
+//
+// 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=valhall3 %s | FileCheck %s --check-prefixes=VALHALL,VALHALL3
+// 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-g710 %s | FileCheck %s --check-prefixes=VALHALL,VALHALL3
+// 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-g510 %s | FileCheck %s --check-prefixes=VALHALL,VALHALL3
+// 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-g310 %s | FileCheck %s --check-prefixes=VALHALL,VALHALL3
+//
+// 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=valhall4 %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=mali-g715 %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=mali-g615 %s | FileCheck %s --check-prefixes=VALHALL,VALHALL4
+
+// VALHALL1: target_info = #iree_gpu.target<arch = "valhall1",
+// VALHALL2: target_info = #iree_gpu.target<arch = "valhall2",
+// VALHALL3: target_info = #iree_gpu.target<arch = "valhall3",
+// VALHALL4: target_info = #iree_gpu.target<arch = "valhall4",
+// VALHALL-SAME: features = "spirv:v1.6,cap:Shader"
+// VALHALL-SAME: wgp = <compute =
+// VALHALL-SAME: fp32|fp16|int32|int16|int8
+// VALHALL-SAME: storage =
+// VALHALL-SAME: b64|b32|b16|b8
+// VALHALL-SAME: subgroup =
+// VALHALL-SAME: shuffle|arithmetic
+// VALHALL-SAME: dot =
+// VALHALL-SAME: dp4xi8toi32
+// VALHALL-SAME: subgroup_size_choices = [16]
+// VALHALL-SAME: max_workgroup_sizes = [512, 512, 512]
+// VALHALL-SAME: max_thread_count_per_workgroup = 512
+// VALHALL-SAME: max_workgroup_memory_bytes = 32768
+// VALHALL-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()
+ stream.return %x, %y, %z : index, index, index
+ }
+ builtin.module {
+ func.func @reduce_dispatch() {
+ return
+ }
+ }
+}
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 e2c5103..a3b3bfd 100644
--- a/compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/KnownTargets.cpp
+++ b/compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/KnownTargets.cpp
@@ -958,7 +958,7 @@
return llvm::StringSwitch<StringRef>(target.lower())
.Cases({"mali-g715", "mali-g615"}, "valhall4")
.Cases({"mali-g710", "mali-g510", "mali-g310"}, "valhall3")
- .Case("mali-78", "valhall2")
+ .Case("mali-g78", "valhall2")
.Cases({"mali-g77", "mali-g57"}, "valhall1")
.Default("");
}
