Drop lists of VulkanSPIRV flags. (#16680)
Previously this could take a list of `--iree-vulkan-target-triple` and
`--iree-vulkan-target-env` flags. Other backends only support a single
flag and as we move towards multi-device, we will be standardizing on
device spec files with `#hal.device.target`s and their executable
formats, _not_ loose flags for each backend.
Now only one of those flags can be set at a time and both are a single
value, not a list.
Progress on https://github.com/openxla/iree/issues/15468 (the flag
parsing code in here was not compatible with how compiler plugins are
currently configured).
diff --git a/compiler/src/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/VulkanSPIRVTarget.cpp b/compiler/src/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/VulkanSPIRVTarget.cpp
index 11aff8e..015a8b3 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/VulkanSPIRVTarget.cpp
+++ b/compiler/src/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/VulkanSPIRVTarget.cpp
@@ -37,17 +37,15 @@
// llvm::cl::OptionCategory halVulkanSPIRVOptionsCategory(
// "IREE Vulkan/SPIR-V backend options");
- static llvm::cl::list<std::string> clVulkanTargetTriples{
+ static llvm::cl::opt<std::string> clVulkanTargetTriple(
"iree-vulkan-target-triple",
llvm::cl::desc(
- "Vulkan target triple controlling the SPIR-V environment."),
- };
+ "Vulkan target triple controlling the SPIR-V environment."));
- static llvm::cl::list<std::string> clVulkanTargetEnvs{
+ static llvm::cl::opt<std::string> clVulkanTargetEnv(
"iree-vulkan-target-env",
llvm::cl::desc(
- "Vulkan target environment as #vk.target_env attribute assembly."),
- };
+ "Vulkan target environment as #vk.target_env attribute assembly."));
static llvm::cl::opt<bool> clVulkanIndirectBindings(
"iree-vulkan-experimental-indirect-bindings",
@@ -55,33 +53,14 @@
llvm::cl::init(false));
VulkanSPIRVTargetOptions targetOptions;
-
- int tripleCount = clVulkanTargetTriples.getNumOccurrences();
- int envCount = clVulkanTargetEnvs.getNumOccurrences();
- int tripleIdx = 0;
- int envIdx = 0;
-
- // Get a flat list of target triples and environments following the original
- // order specified via the command line.
- SmallVector<std::string> vulkanTargetTriplesAndEnvs;
- for (int i = 0, e = tripleCount + envCount; i < e; ++i) {
- if (tripleIdx >= tripleCount) {
- vulkanTargetTriplesAndEnvs.push_back(clVulkanTargetEnvs[envIdx++]);
- continue;
- }
- if (envIdx >= envCount) {
- vulkanTargetTriplesAndEnvs.push_back(clVulkanTargetTriples[tripleIdx++]);
- continue;
- }
- if (clVulkanTargetTriples.getPosition(tripleIdx) >
- clVulkanTargetEnvs.getPosition(envIdx)) {
- vulkanTargetTriplesAndEnvs.push_back(clVulkanTargetEnvs[envIdx++]);
- } else {
- vulkanTargetTriplesAndEnvs.push_back(clVulkanTargetTriples[tripleIdx++]);
- }
+ if (!clVulkanTargetEnv.empty()) {
+ // TODO(scotttodd): assert if triple is set too? (mutually exclusive flags)
+ targetOptions.targetTripleOrEnv = clVulkanTargetEnv;
+ } else if (!clVulkanTargetTriple.empty()) {
+ targetOptions.targetTripleOrEnv = clVulkanTargetTriple;
+ } else {
+ targetOptions.targetTripleOrEnv = "unknown-unknown-unknown";
}
- targetOptions.targetTriplesAndEnvs = vulkanTargetTriplesAndEnvs;
-
targetOptions.indirectBindings = clVulkanIndirectBindings;
return targetOptions;
@@ -93,10 +72,12 @@
MLIRContext *context) {
if (!vulkanTargetTripleOrEnv.empty()) {
if (vulkanTargetTripleOrEnv[0] != '#') {
+ // Parse target triple.
return convertTargetEnv(
Vulkan::getTargetEnvForTriple(context, vulkanTargetTripleOrEnv));
}
+ // Parse `#vk.target_env<...` attribute assembly.
if (auto attr = parseAttribute(vulkanTargetTripleOrEnv, context)) {
if (auto vkTargetEnv = llvm::dyn_cast<Vulkan::TargetEnvAttr>(attr)) {
return convertTargetEnv(vkTargetEnv);
@@ -148,19 +129,9 @@
MLIRContext *context, StringRef deviceID, DictionaryAttr deviceConfigAttr,
SmallVectorImpl<IREE::HAL::ExecutableTargetAttr> &executableTargetAttrs)
const override {
- // Select SPIR-V environments to compile for.
- for (std::string targetTripleOrEnv : options_.targetTriplesAndEnvs) {
- executableTargetAttrs.push_back(getExecutableTarget(
- context, getSPIRVTargetEnv(targetTripleOrEnv, context),
- options_.indirectBindings));
- }
-
- // If no environment specified, populate with a minimal target.
- if (executableTargetAttrs.empty()) {
- executableTargetAttrs.push_back(getExecutableTarget(
- context, getSPIRVTargetEnv("unknown-unknown-unknown", context),
- options_.indirectBindings));
- }
+ executableTargetAttrs.push_back(getExecutableTarget(
+ context, getSPIRVTargetEnv(options_.targetTripleOrEnv, context),
+ options_.indirectBindings));
}
IREE::HAL::ExecutableTargetAttr
diff --git a/compiler/src/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/VulkanSPIRVTarget.h b/compiler/src/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/VulkanSPIRVTarget.h
index 5365563..6175f1c 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/VulkanSPIRVTarget.h
+++ b/compiler/src/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/VulkanSPIRVTarget.h
@@ -10,15 +10,13 @@
#include <functional>
#include <string>
-#include "llvm/ADT/SmallVector.h"
-
namespace mlir::iree_compiler::IREE::HAL {
// Options controlling the SPIR-V translation.
struct VulkanSPIRVTargetOptions {
- // Vulkan target environments, either as #vk.target_env attribute assembly
+ // Vulkan target environment, either as #vk.target_env attribute assembly
// or as a Vulkan target triple.
- llvm::SmallVector<std::string> targetTriplesAndEnvs;
+ std::string targetTripleOrEnv;
// Whether to use indirect bindings for all generated dispatches.
bool indirectBindings = false;
};
diff --git a/compiler/src/iree/compiler/Dialect/Vulkan/Utils/test/BUILD.bazel b/compiler/src/iree/compiler/Dialect/Vulkan/Utils/test/BUILD.bazel
index 084c2b1..687fa49 100644
--- a/compiler/src/iree/compiler/Dialect/Vulkan/Utils/test/BUILD.bazel
+++ b/compiler/src/iree/compiler/Dialect/Vulkan/Utils/test/BUILD.bazel
@@ -25,7 +25,6 @@
name = "lit",
srcs = enforce_glob(
[
- "multiple_target_env_conversion.mlir",
"target_env_conversion.mlir",
],
include = ["*.mlir"],
diff --git a/compiler/src/iree/compiler/Dialect/Vulkan/Utils/test/CMakeLists.txt b/compiler/src/iree/compiler/Dialect/Vulkan/Utils/test/CMakeLists.txt
index 524a0ba..bb5cbe5 100644
--- a/compiler/src/iree/compiler/Dialect/Vulkan/Utils/test/CMakeLists.txt
+++ b/compiler/src/iree/compiler/Dialect/Vulkan/Utils/test/CMakeLists.txt
@@ -18,7 +18,6 @@
NAME
lit
SRCS
- "multiple_target_env_conversion.mlir"
"target_env_conversion.mlir"
TOOLS
FileCheck
diff --git a/compiler/src/iree/compiler/Dialect/Vulkan/Utils/test/multiple_target_env_conversion.mlir b/compiler/src/iree/compiler/Dialect/Vulkan/Utils/test/multiple_target_env_conversion.mlir
deleted file mode 100644
index e51577f..0000000
--- a/compiler/src/iree/compiler/Dialect/Vulkan/Utils/test/multiple_target_env_conversion.mlir
+++ /dev/null
@@ -1,44 +0,0 @@
-// RUN: iree-opt --pass-pipeline='builtin.module(iree-hal-assign-target-devices{targetBackends=vulkan-spirv},iree-hal-transformation-pipeline{serialize-executables=false})' \
-// RUN: --iree-vulkan-target-triple=rdna3-7900xtx-windows \
-// RUN: --iree-vulkan-target-env="#vk.target_env<v1.1, r(120), [VK_KHR_spirv_1_4, VK_KHR_storage_buffer_storage_class], AMD:DiscreteGPU, #vk.caps<maxComputeSharedMemorySize = 16384, maxComputeWorkGroupInvocations = 1024, maxComputeWorkGroupSize = dense<[128, 8, 4]>: vector<3xi32>, subgroupFeatures = 63 : i32, subgroupSize = 4 >>" \
-// RUN: --iree-vulkan-target-triple=valhall-unknown-android31 \
-// RUN: %s | FileCheck %s
-
-// CHECK: #[[RDNA3:.+]] = {{.*}} #spirv.target_env<#spirv.vce<v1.6, [Shader, Float64, Float16, Int64, Int16, Int8, StorageBuffer16BitAccess, StorageUniform16, StoragePushConstant16, StorageBuffer8BitAccess, UniformAndStorageBuffer8BitAccess, StoragePushConstant8, GroupNonUniform, GroupNonUniformVote, GroupNonUniformArithmetic, GroupNonUniformBallot, GroupNonUniformShuffle, GroupNonUniformShuffleRelative, GroupNonUniformClustered, GroupNonUniformQuad, PhysicalStorageBufferAddresses, VariablePointers, VariablePointersStorageBuffer, DotProduct, DotProductInputAll, DotProductInput4x8BitPacked, DotProductInput4x8Bit, CooperativeMatrixKHR],
-// CHECK-SAME: [SPV_KHR_16bit_storage, SPV_KHR_8bit_storage, SPV_KHR_integer_dot_product, SPV_KHR_storage_buffer_storage_class, SPV_KHR_physical_storage_buffer, SPV_KHR_variable_pointers, SPV_KHR_cooperative_matrix]>,
-// CHECK-SAME: api=Vulkan, AMD:DiscreteGPU, #spirv.resource_limits<max_compute_shared_memory_size = 65536, max_compute_workgroup_invocations = 1024, max_compute_workgroup_size = [1024, 1024, 1024], subgroup_size = 64, min_subgroup_size = 32, max_subgroup_size = 64, cooperative_matrix_properties_khr = [#spirv.coop_matrix_props_khr<m_size = 16, n_size = 16, k_size = 16, a_type = i8, b_type = i8, c_type = i32, result_type = i32, acc_sat = false, scope = <Subgroup>>, #spirv.coop_matrix_props_khr<m_size = 16, n_size = 16, k_size = 16, a_type = f16, b_type = f16, c_type = f16, result_type = f16, acc_sat = false, scope = <Subgroup>>, #spirv.coop_matrix_props_khr<m_size = 16, n_size = 16, k_size = 16, a_type = f16, b_type = f16, c_type = f32, result_type = f32, acc_sat = false, scope = <Subgroup>>]>>
-
-// CHECK: #[[ENV:.+]] = {{.*}} #spirv.target_env<#spirv.vce<v1.4, [Shader, GroupNonUniform, GroupNonUniformVote, GroupNonUniformArithmetic, GroupNonUniformBallot, GroupNonUniformShuffle, GroupNonUniformShuffleRelative],
-// CHECK-SAME: [SPV_KHR_storage_buffer_storage_class]>,
-// CHECK-SAME: api=Vulkan, AMD:DiscreteGPU, #spirv.resource_limits<max_compute_workgroup_invocations = 1024, max_compute_workgroup_size = [128, 8, 4], subgroup_size = 4, cooperative_matrix_properties_khr = []>>
-
-// CHECK: #[[VALHALL:.+]] = {{.*}} #spirv.target_env<#spirv.vce<v1.4, [Shader, Float16, Int16, Int8, StorageBuffer16BitAccess, StorageUniform16, StoragePushConstant16, StorageBuffer8BitAccess, UniformAndStorageBuffer8BitAccess, StoragePushConstant8, GroupNonUniform, GroupNonUniformVote, GroupNonUniformArithmetic, GroupNonUniformBallot, GroupNonUniformShuffle, GroupNonUniformShuffleRelative, GroupNonUniformClustered, GroupNonUniformQuad, VariablePointers, VariablePointersStorageBuffer, DotProduct, DotProductInputAll, DotProductInput4x8BitPacked, DotProductInput4x8Bit],
-// CHECK-SAME: [SPV_KHR_16bit_storage, SPV_KHR_8bit_storage, SPV_KHR_integer_dot_product, SPV_KHR_storage_buffer_storage_class, SPV_KHR_variable_pointers]>,
-// CHECK-SAME: api=Vulkan, ARM:IntegratedGPU, #spirv.resource_limits<max_compute_shared_memory_size = 32768, max_compute_workgroup_invocations = 512, max_compute_workgroup_size = [512, 512, 512], subgroup_size = 16, cooperative_matrix_properties_khr = []>>
-
-// Verify that the order of target environments matches what the user specified.
-
-// CHECK: [#[[RDNA3]], #[[ENV]], #[[VALHALL]]]
-
-stream.executable public @reduce_dispatch {
- stream.executable.export @reduce_dispatch workgroups(%arg0: index) -> (index, index, index) {
- %x, %y, %z = flow.dispatch.workgroup_count_from_dag_root %arg0
- stream.return %x, %y, %z : index, index, index
- }
- builtin.module {
- func.func @reduce_dispatch(%arg0_binding: !stream.binding, %arg1_binding: !stream.binding) {
- %c0 = arith.constant 0 : index
- %arg0 = stream.binding.subspan %arg0_binding[%c0] : !stream.binding -> !flow.dispatch.tensor<readonly:tensor<16xf32>>
- %arg1 = stream.binding.subspan %arg1_binding[%c0] : !stream.binding -> !flow.dispatch.tensor<writeonly:tensor<f32>>
- %0 = tensor.empty() : tensor<f32>
- %1 = flow.dispatch.tensor.load %arg0, offsets=[0], sizes=[16], strides=[1] : !flow.dispatch.tensor<readonly:tensor<16xf32>> -> tensor<16xf32>
- %3 = linalg.generic {indexing_maps = [affine_map<(d0) -> (d0)>, affine_map<(d0) -> ()>], iterator_types = ["reduction"]} ins(%1 : tensor<16xf32>) outs(%0 : tensor<f32>) {
- ^bb0(%arg2: f32, %arg3: f32):
- %4 = arith.addf %arg2, %arg3 : f32
- linalg.yield %4 : f32
- } -> tensor<f32>
- flow.dispatch.tensor.store %3, %arg1, offsets=[], sizes=[], strides=[] : tensor<f32> -> !flow.dispatch.tensor<writeonly:tensor<f32>>
- return
- }
- }
-}
diff --git a/compiler/src/iree/compiler/Dialect/Vulkan/Utils/test/target_env_conversion.mlir b/compiler/src/iree/compiler/Dialect/Vulkan/Utils/test/target_env_conversion.mlir
index 69d1010..3a23031 100644
--- a/compiler/src/iree/compiler/Dialect/Vulkan/Utils/test/target_env_conversion.mlir
+++ b/compiler/src/iree/compiler/Dialect/Vulkan/Utils/test/target_env_conversion.mlir
@@ -7,6 +7,7 @@
// RUN: iree-opt --pass-pipeline='builtin.module(iree-hal-assign-target-devices{targetBackends=vulkan-spirv},iree-hal-transformation-pipeline{serialize-executables=false})' --iree-vulkan-target-triple=m1-moltenvk-macos %s | FileCheck %s --check-prefix=M1
// RUN: iree-opt --pass-pipeline='builtin.module(iree-hal-assign-target-devices{targetBackends=vulkan-spirv},iree-hal-transformation-pipeline{serialize-executables=false})' --iree-vulkan-target-triple=arc-770-windows %s | FileCheck %s --check-prefix=ARC
// RUN: iree-opt --pass-pipeline='builtin.module(iree-hal-assign-target-devices{targetBackends=vulkan-spirv},iree-hal-transformation-pipeline{serialize-executables=false})' --iree-vulkan-target-triple=pascal-1080-windows %s | FileCheck %s --check-prefix=PASCAL
+// RUN: iree-opt --pass-pipeline='builtin.module(iree-hal-assign-target-devices{targetBackends=vulkan-spirv},iree-hal-transformation-pipeline{serialize-executables=false})' --iree-vulkan-target-env="#vk.target_env<v1.1, r(120), [VK_KHR_spirv_1_4, VK_KHR_storage_buffer_storage_class], AMD:DiscreteGPU, #vk.caps<maxComputeSharedMemorySize = 16384, maxComputeWorkGroupInvocations = 1024, maxComputeWorkGroupSize = dense<[128, 8, 4]>: vector<3xi32>, subgroupFeatures = 63 : i32, subgroupSize = 4 >>" %s | FileCheck %s --check-prefix=ENV
// TODO(antiagainst): Passing in lenghty strings as command-line options is not
// optimal. We should consider creating a dedicated test pass to pick up
@@ -56,6 +57,11 @@
// PASCAL-SAME: [SPV_KHR_16bit_storage, SPV_KHR_8bit_storage, SPV_KHR_integer_dot_product, SPV_KHR_storage_buffer_storage_class, SPV_KHR_physical_storage_buffer, SPV_KHR_variable_pointers, SPV_KHR_cooperative_matrix]>,
// PASCAL-SAME: api=Vulkan, NVIDIA:DiscreteGPU, #spirv.resource_limits<max_compute_shared_memory_size = 49152, max_compute_workgroup_invocations = 1536, max_compute_workgroup_size = [1536, 1024, 64], min_subgroup_size = 32, max_subgroup_size = 32, cooperative_matrix_properties_khr = []>>}>
+// ENV: #spirv.target_env<#spirv.vce<v1.4,
+// ENV-SAME: [Shader, GroupNonUniform, GroupNonUniformVote, GroupNonUniformArithmetic, GroupNonUniformBallot, GroupNonUniformShuffle, GroupNonUniformShuffleRelative],
+// ENV-SAME: [SPV_KHR_storage_buffer_storage_class]>,
+// ENV-SAME: api=Vulkan, AMD:DiscreteGPU, #spirv.resource_limits<max_compute_workgroup_invocations = 1024, max_compute_workgroup_size = [128, 8, 4], subgroup_size = 4, cooperative_matrix_properties_khr = []>>
+
stream.executable public @reduce_dispatch {
stream.executable.export @reduce_dispatch workgroups(%arg0: index) -> (index, index, index) {
%x, %y, %z = flow.dispatch.workgroup_count_from_dag_root %arg0
