Typing IREE::HAL::DeviceTargetAttr executable targets. (#16588)
This allows the device configuration to be independent from the
executable target attributes which is useful for devices that may share
the same configuration but different executable targets.
diff --git a/compiler/plugins/target/CUDA/CUDATarget.cpp b/compiler/plugins/target/CUDA/CUDATarget.cpp
index b5ef37f..6381319 100644
--- a/compiler/plugins/target/CUDA/CUDATarget.cpp
+++ b/compiler/plugins/target/CUDA/CUDATarget.cpp
@@ -384,12 +384,17 @@
Builder b(context);
SmallVector<NamedAttribute> configItems;
- configItems.emplace_back(b.getStringAttr("executable_targets"),
- getExecutableTargets(context));
+ // TODO: device configuration attrs.
auto configAttr = b.getDictionaryAttr(configItems);
+
+ // If we had multiple target environments we would generate one target attr
+ // per environment, with each setting its own environment attribute.
+ SmallVector<IREE::HAL::ExecutableTargetAttr> targetAttrs;
+ targetAttrs.push_back(getExecutableTarget(context));
+
return IREE::HAL::DeviceTargetAttr::get(
- context, b.getStringAttr(deviceID()), configAttr);
+ context, b.getStringAttr(deviceID()), configAttr, targetAttrs);
}
void buildConfigurationPassPipeline(IREE::HAL::ExecutableVariantOp variantOp,
@@ -640,14 +645,6 @@
}
private:
- ArrayAttr getExecutableTargets(MLIRContext *context) const {
- SmallVector<Attribute> targetAttrs;
- // If we had multiple target environments we would generate one target attr
- // per environment, with each setting its own environment attribute.
- targetAttrs.push_back(getExecutableTarget(context));
- return ArrayAttr::get(context, targetAttrs);
- }
-
IREE::HAL::ExecutableTargetAttr
getExecutableTarget(MLIRContext *context) const {
Builder b(context);
diff --git a/compiler/plugins/target/CUDA/test/smoketest.mlir b/compiler/plugins/target/CUDA/test/smoketest.mlir
index e3ac33b..56285fa 100644
--- a/compiler/plugins/target/CUDA/test/smoketest.mlir
+++ b/compiler/plugins/target/CUDA/test/smoketest.mlir
@@ -5,11 +5,9 @@
module attributes {
hal.device.targets = [
- #hal.device.target<"cuda", {
- executable_targets = [
- #hal.executable.target<"cuda", "cuda-nvptx-fb">
- ]
- }>
+ #hal.device.target<"cuda", [
+ #hal.executable.target<"cuda", "cuda-nvptx-fb">
+ ]>
]
} {
diff --git a/compiler/plugins/target/MetalSPIRV/MetalSPIRVTarget.cpp b/compiler/plugins/target/MetalSPIRV/MetalSPIRVTarget.cpp
index 28269e5..9ea8853 100644
--- a/compiler/plugins/target/MetalSPIRV/MetalSPIRVTarget.cpp
+++ b/compiler/plugins/target/MetalSPIRV/MetalSPIRVTarget.cpp
@@ -123,12 +123,16 @@
Builder b(context);
SmallVector<NamedAttribute> configItems;
- configItems.emplace_back(b.getStringAttr("executable_targets"),
- getExecutableTargets(context));
-
auto configAttr = b.getDictionaryAttr(configItems);
+
+ // If we had multiple target environments we would generate one target attr
+ // per environment, with each setting its own environment attribute.
+ SmallVector<IREE::HAL::ExecutableTargetAttr> targetAttrs;
+ targetAttrs.push_back(
+ getExecutableTarget(context, getMetalTargetEnv(context)));
+
return IREE::HAL::DeviceTargetAttr::get(
- context, b.getStringAttr(deviceID()), configAttr);
+ context, b.getStringAttr(deviceID()), configAttr, targetAttrs);
}
void buildConfigurationPassPipeline(IREE::HAL::ExecutableVariantOp variantOp,
@@ -282,15 +286,6 @@
}
private:
- ArrayAttr getExecutableTargets(MLIRContext *context) const {
- SmallVector<Attribute> targetAttrs;
- // If we had multiple target environments we would generate one target attr
- // per environment, with each setting its own environment attribute.
- targetAttrs.push_back(
- getExecutableTarget(context, getMetalTargetEnv(context)));
- return ArrayAttr::get(context, targetAttrs);
- }
-
IREE::HAL::ExecutableTargetAttr
getExecutableTarget(MLIRContext *context,
spirv::TargetEnvAttr targetEnv) const {
diff --git a/compiler/plugins/target/MetalSPIRV/test/smoketest.mlir b/compiler/plugins/target/MetalSPIRV/test/smoketest.mlir
index 9a39573..84dc61e 100644
--- a/compiler/plugins/target/MetalSPIRV/test/smoketest.mlir
+++ b/compiler/plugins/target/MetalSPIRV/test/smoketest.mlir
@@ -2,13 +2,11 @@
module attributes {
hal.device.targets = [
- #hal.device.target<"metal", {
- executable_targets = [
- #hal.executable.target<"metal-spirv", "metal-msl-fb", {
- spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spirv.resource_limits<>>
- }>
- ]
- }>
+ #hal.device.target<"metal", [
+ #hal.executable.target<"metal-spirv", "metal-msl-fb", {
+ spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spirv.resource_limits<>>
+ }>
+ ]>
]
} {
diff --git a/compiler/plugins/target/ROCM/ROCMTarget.cpp b/compiler/plugins/target/ROCM/ROCMTarget.cpp
index 21406b4..25d1eee 100644
--- a/compiler/plugins/target/ROCM/ROCMTarget.cpp
+++ b/compiler/plugins/target/ROCM/ROCMTarget.cpp
@@ -157,12 +157,15 @@
// synchronous mode.
configItems.emplace_back(b.getStringAttr("legacy_sync"), b.getUnitAttr());
- configItems.emplace_back(b.getStringAttr("executable_targets"),
- getExecutableTargets(context));
-
auto configAttr = b.getDictionaryAttr(configItems);
+
+ // If we had multiple target environments we would generate one target attr
+ // per environment, with each setting its own environment attribute.
+ SmallVector<IREE::HAL::ExecutableTargetAttr> targetAttrs;
+ targetAttrs.push_back(getExecutableTarget(context));
+
return IREE::HAL::DeviceTargetAttr::get(
- context, b.getStringAttr(deviceID()), configAttr);
+ context, b.getStringAttr(deviceID()), configAttr, targetAttrs);
}
// Performs optimizations on |module| (including LTO-style whole-program
// ones). Inspired by code section in
@@ -457,14 +460,6 @@
}
private:
- ArrayAttr getExecutableTargets(MLIRContext *context) const {
- SmallVector<Attribute> targetAttrs;
- // If we had multiple target environments we would generate one target attr
- // per environment, with each setting its own environment attribute.
- targetAttrs.push_back(getExecutableTarget(context));
- return ArrayAttr::get(context, targetAttrs);
- }
-
IREE::HAL::ExecutableTargetAttr
getExecutableTarget(MLIRContext *context) const {
Builder b(context);
diff --git a/compiler/plugins/target/ROCM/test/smoketest.mlir b/compiler/plugins/target/ROCM/test/smoketest.mlir
index 201e123..2a46244 100644
--- a/compiler/plugins/target/ROCM/test/smoketest.mlir
+++ b/compiler/plugins/target/ROCM/test/smoketest.mlir
@@ -2,11 +2,9 @@
module attributes {
hal.device.targets = [
- #hal.device.target<"rocm", {
- executable_targets = [
- #hal.executable.target<"rocm", "rocm-hsaco-fb">
- ]
- }>
+ #hal.device.target<"rocm", [
+ #hal.executable.target<"rocm", "rocm-hsaco-fb">
+ ]>
]
} {
diff --git a/compiler/plugins/target/VMVX/VMVXTarget.cpp b/compiler/plugins/target/VMVX/VMVXTarget.cpp
index bd5eae4..abf72c3 100644
--- a/compiler/plugins/target/VMVX/VMVXTarget.cpp
+++ b/compiler/plugins/target/VMVX/VMVXTarget.cpp
@@ -67,12 +67,16 @@
Builder b(context);
SmallVector<NamedAttribute> configItems;
- configItems.emplace_back(b.getStringAttr("executable_targets"),
- getExecutableTargets(context));
-
auto configAttr = b.getDictionaryAttr(configItems);
+
+ // If we had multiple target environments we would generate one target attr
+ // per environment, with each setting its own environment attribute.
+ SmallVector<IREE::HAL::ExecutableTargetAttr> targetAttrs;
+ targetAttrs.push_back(getVMVXExecutableTarget(
+ options.enableMicrokernels, context, "vmvx", "vmvx-bytecode-fb"));
+
return IREE::HAL::DeviceTargetAttr::get(
- context, b.getStringAttr(deviceID()), configAttr);
+ context, b.getStringAttr(deviceID()), configAttr, targetAttrs);
}
std::optional<IREE::HAL::DeviceTargetAttr>
@@ -164,14 +168,6 @@
}
private:
- ArrayAttr getExecutableTargets(MLIRContext *context) const {
- SmallVector<Attribute> targetAttrs;
- // This is where we would multiversion.
- targetAttrs.push_back(getVMVXExecutableTarget(
- options.enableMicrokernels, context, "vmvx", "vmvx-bytecode-fb"));
- return ArrayAttr::get(context, targetAttrs);
- }
-
const VMVXOptions &options;
};
@@ -191,12 +187,16 @@
Builder b(context);
SmallVector<NamedAttribute> configItems;
- configItems.emplace_back(b.getStringAttr("executable_targets"),
- getExecutableTargets(context));
-
auto configAttr = b.getDictionaryAttr(configItems);
+
+ // If we had multiple target environments we would generate one target attr
+ // per environment, with each setting its own environment attribute.
+ SmallVector<IREE::HAL::ExecutableTargetAttr> targetAttrs;
+ targetAttrs.push_back(getVMVXExecutableTarget(
+ options.enableMicrokernels, context, "vmvx-inline", "vmvx-ir"));
+
return IREE::HAL::DeviceTargetAttr::get(
- context, b.getStringAttr(deviceID()), configAttr);
+ context, b.getStringAttr(deviceID()), configAttr, targetAttrs);
}
void buildConfigurationPassPipeline(IREE::HAL::ExecutableVariantOp variantOp,
@@ -210,14 +210,6 @@
}
private:
- ArrayAttr getExecutableTargets(MLIRContext *context) const {
- SmallVector<Attribute> targetAttrs;
- // This is where we would multiversion.
- targetAttrs.push_back(getVMVXExecutableTarget(
- options.enableMicrokernels, context, "vmvx-inline", "vmvx-ir"));
- return ArrayAttr::get(context, targetAttrs);
- }
-
const VMVXOptions &options;
};
diff --git a/compiler/plugins/target/VMVX/test/smoketest.mlir b/compiler/plugins/target/VMVX/test/smoketest.mlir
index 7a67a80..ef56028 100644
--- a/compiler/plugins/target/VMVX/test/smoketest.mlir
+++ b/compiler/plugins/target/VMVX/test/smoketest.mlir
@@ -2,11 +2,9 @@
module attributes {
hal.device.targets = [
- #hal.device.target<"vmvx", {
- executable_targets = [
- #hal.executable.target<"vmvx", "vmvx-bytecode-fb">
- ]
- }>
+ #hal.device.target<"vmvx", [
+ #hal.executable.target<"vmvx", "vmvx-bytecode-fb">
+ ]>
]
} {
diff --git a/compiler/plugins/target/WebGPUSPIRV/WebGPUSPIRVTarget.cpp b/compiler/plugins/target/WebGPUSPIRV/WebGPUSPIRVTarget.cpp
index 1f57ad5..dd39ebd 100644
--- a/compiler/plugins/target/WebGPUSPIRV/WebGPUSPIRVTarget.cpp
+++ b/compiler/plugins/target/WebGPUSPIRV/WebGPUSPIRVTarget.cpp
@@ -76,12 +76,16 @@
Builder b(context);
SmallVector<NamedAttribute> configItems;
- configItems.emplace_back(b.getStringAttr("executable_targets"),
- getExecutableTargets(context));
-
auto configAttr = b.getDictionaryAttr(configItems);
- return IREE::HAL::DeviceTargetAttr::get(
- context, b.getStringAttr(deviceID()), configAttr);
+
+ // If we had multiple target environments we would generate one target attr
+ // per environment, with each setting its own environment attribute.
+ SmallVector<IREE::HAL::ExecutableTargetAttr> targetAttrs;
+ targetAttrs.push_back(
+ getExecutableTarget(context, getWebGPUTargetEnv(context)));
+
+ return IREE::HAL::DeviceTargetAttr::get(context, b.getStringAttr("webgpu"),
+ configAttr, targetAttrs);
}
void buildConfigurationPassPipeline(IREE::HAL::ExecutableVariantOp variantOp,
@@ -255,15 +259,6 @@
}
private:
- ArrayAttr getExecutableTargets(MLIRContext *context) const {
- SmallVector<Attribute> targetAttrs;
- // If we had multiple target environments we would generate one target attr
- // per environment, with each setting its own environment attribute.
- targetAttrs.push_back(
- getExecutableTarget(context, getWebGPUTargetEnv(context)));
- return ArrayAttr::get(context, targetAttrs);
- }
-
IREE::HAL::ExecutableTargetAttr
getExecutableTarget(MLIRContext *context,
spirv::TargetEnvAttr targetEnv) const {
diff --git a/compiler/plugins/target/WebGPUSPIRV/test/smoketest.mlir b/compiler/plugins/target/WebGPUSPIRV/test/smoketest.mlir
index 3613d79..1a17240 100644
--- a/compiler/plugins/target/WebGPUSPIRV/test/smoketest.mlir
+++ b/compiler/plugins/target/WebGPUSPIRV/test/smoketest.mlir
@@ -3,13 +3,11 @@
#map = affine_map<(d0) -> (d0)>
module attributes {
hal.device.targets = [
- #hal.device.target<"webgpu", {
- executable_targets = [
- #hal.executable.target<"webgpu-spirv", "webgpu-wgsl-fb", {
- spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spirv.resource_limits<>>
- }>
- ]
- }>
+ #hal.device.target<"webgpu", [
+ #hal.executable.target<"webgpu-spirv", "webgpu-wgsl-fb", {
+ spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spirv.resource_limits<>>
+ }>
+ ]>
]
} {
diff --git a/compiler/src/iree/compiler/Codegen/LLVMCPU/test/transform_dialect_bufferize.mlir b/compiler/src/iree/compiler/Codegen/LLVMCPU/test/transform_dialect_bufferize.mlir
index 466c259..2adc175 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMCPU/test/transform_dialect_bufferize.mlir
+++ b/compiler/src/iree/compiler/Codegen/LLVMCPU/test/transform_dialect_bufferize.mlir
@@ -1,8 +1,8 @@
// RUN: iree-opt %s -iree-transform-dialect-interpreter -transform-dialect-drop-schedule | FileCheck %s
-#device_target_cpu = #hal.device.target<"llvm-cpu", {executable_targets = [#hal.executable.target<"llvm-cpu", "embedded-elf-x86_64", {cpu_features = "", data_layout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-f80:128-n8:16:32:64-S128", native_vector_size = 16 : index, target_triple = "x86_64-none-elf"}>]}>
-#pipeline_layout = #hal.pipeline.layout<push_constants = 0, sets = [#hal.descriptor_set.layout<0, bindings = [#hal.descriptor_set.binding<0, storage_buffer>, #hal.descriptor_set.binding<1, storage_buffer>, #hal.descriptor_set.binding<2, storage_buffer>]>]>
#executable_target_embedded_elf_x86_64_ = #hal.executable.target<"llvm-cpu", "embedded-elf-x86_64", {cpu_features = "", data_layout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-f80:128-n8:16:32:64-S128", native_vector_size = 16 : index, target_triple = "x86_64-none-elf"}>
+#device_target_cpu = #hal.device.target<"llvm-cpu", [#executable_target_embedded_elf_x86_64_]>
+#pipeline_layout = #hal.pipeline.layout<push_constants = 0, sets = [#hal.descriptor_set.layout<0, bindings = [#hal.descriptor_set.binding<0, storage_buffer>, #hal.descriptor_set.binding<1, storage_buffer>, #hal.descriptor_set.binding<2, storage_buffer>]>]>
hal.executable private @pad_matmul_static_dispatch_0 {
hal.executable.variant public @embedded_elf_x86_64 target(#executable_target_embedded_elf_x86_64_) {
diff --git a/compiler/src/iree/compiler/Dialect/HAL/Conversion/StreamToHAL/test/cmd_ops.mlir b/compiler/src/iree/compiler/Dialect/HAL/Conversion/StreamToHAL/test/cmd_ops.mlir
index 9093f31..7cdd991 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/Conversion/StreamToHAL/test/cmd_ops.mlir
+++ b/compiler/src/iree/compiler/Dialect/HAL/Conversion/StreamToHAL/test/cmd_ops.mlir
@@ -166,9 +166,10 @@
#executable_target_aarch64 = #hal.executable.target<"llvm-cpu", "embedded-elf-aarch64">
#executable_target_x86_64 = #hal.executable.target<"llvm-cpu", "embedded-elf-x86_64">
-#device_target_cpu = #hal.device.target<"llvm-cpu", {
- executable_targets = [#executable_target_aarch64, #executable_target_x86_64]
-}>
+#device_target_cpu = #hal.device.target<"llvm-cpu", [
+ #executable_target_aarch64,
+ #executable_target_x86_64
+]>
#pipeline_layout = #hal.pipeline.layout<push_constants = 0, sets = [
#hal.descriptor_set.layout<0, bindings = [
#hal.descriptor_set.binding<4, storage_buffer>
diff --git a/compiler/src/iree/compiler/Dialect/HAL/IR/HALAttrs.cpp b/compiler/src/iree/compiler/Dialect/HAL/IR/HALAttrs.cpp
index 8ad4f14..82d38ae 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/IR/HALAttrs.cpp
+++ b/compiler/src/iree/compiler/Dialect/HAL/IR/HALAttrs.cpp
@@ -93,27 +93,54 @@
StringRef deviceID) {
// TODO(benvanik): query default configuration from the target backend.
return get(context, StringAttr::get(context, deviceID),
- DictionaryAttr::get(context));
+ DictionaryAttr::get(context), {});
}
// static
Attribute DeviceTargetAttr::parse(AsmParser &p, Type type) {
StringAttr deviceIDAttr;
DictionaryAttr configAttr;
+ SmallVector<IREE::HAL::ExecutableTargetAttr> executableTargetAttrs;
// `<"device-id"`
if (failed(p.parseLess()) || failed(p.parseAttribute(deviceIDAttr))) {
return {};
}
- // `, {config}`
- if (succeeded(p.parseOptionalComma()) &&
- failed(p.parseAttribute(configAttr))) {
- return {};
+ // `, `
+ if (succeeded(p.parseOptionalComma())) {
+ if (succeeded(p.parseOptionalLSquare())) {
+ // `[targets, ...]` (optional)
+ do {
+ IREE::HAL::ExecutableTargetAttr executableTargetAttr;
+ if (failed(p.parseAttribute(executableTargetAttr)))
+ return {};
+ executableTargetAttrs.push_back(executableTargetAttr);
+ } while (succeeded(p.parseOptionalComma()));
+ if (failed(p.parseRSquare()))
+ return {};
+ } else {
+ // `{config dict}` (optional)
+ if (failed(p.parseAttribute(configAttr)))
+ return {};
+ // `, [targets, ...]` (optional)
+ if (succeeded(p.parseOptionalComma())) {
+ if (failed(p.parseLSquare()))
+ return {};
+ do {
+ IREE::HAL::ExecutableTargetAttr executableTargetAttr;
+ if (failed(p.parseAttribute(executableTargetAttr)))
+ return {};
+ executableTargetAttrs.push_back(executableTargetAttr);
+ } while (succeeded(p.parseOptionalComma()));
+ if (failed(p.parseRSquare()))
+ return {};
+ }
+ }
}
// `>`
if (failed(p.parseGreater())) {
return {};
}
- return get(p.getContext(), deviceIDAttr, configAttr);
+ return get(p.getContext(), deviceIDAttr, configAttr, executableTargetAttrs);
}
void DeviceTargetAttr::print(AsmPrinter &p) const {
@@ -125,6 +152,15 @@
os << ", ";
p.printAttribute(configAttr);
}
+ auto executableTargetAttrs = getExecutableTargets();
+ if (!executableTargetAttrs.empty()) {
+ os << ", [";
+ llvm::interleaveComma(executableTargetAttrs, os,
+ [&](auto executableTargetAttr) {
+ p.printAttribute(executableTargetAttr);
+ });
+ os << "]";
+ }
os << ">";
}
@@ -137,20 +173,6 @@
return configAttr && configAttr.get(name);
}
-SmallVector<ExecutableTargetAttr, 4> DeviceTargetAttr::getExecutableTargets() {
- SmallVector<ExecutableTargetAttr, 4> resultAttrs;
- auto configAttr = getConfiguration();
- if (configAttr) {
- auto targetsAttr = configAttr.getAs<ArrayAttr>("executable_targets");
- if (targetsAttr) {
- for (auto attr : targetsAttr.getValue()) {
- resultAttrs.push_back(llvm::dyn_cast<ExecutableTargetAttr>(attr));
- }
- }
- }
- return resultAttrs;
-}
-
// static
SmallVector<IREE::HAL::DeviceTargetAttr, 4>
DeviceTargetAttr::lookup(Operation *op) {
diff --git a/compiler/src/iree/compiler/Dialect/HAL/IR/HALAttrs.td b/compiler/src/iree/compiler/Dialect/HAL/IR/HALAttrs.td
index a03879b..80add77 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/IR/HALAttrs.td
+++ b/compiler/src/iree/compiler/Dialect/HAL/IR/HALAttrs.td
@@ -489,16 +489,17 @@
Example:
```mlir
#hal.device.target<"llvm-cpu", {
- executable_targets = [
- #hal.executable.target<"llvm-cpu", "embedded-elf-arm_32">,
- #hal.executable.target<"llvm-cpu", "embedded-elf-arm_64">,
- ]
- }>
+ device_configuration = ...
+ }, [
+ #hal.executable.target<"llvm-cpu", "embedded-elf-arm_32">,
+ #hal.executable.target<"llvm-cpu", "embedded-elf-arm_64">,
+ ]>
```
}];
let parameters = (ins
AttrParameter<"StringAttr", "">:$deviceID,
- AttrParameter<"DictionaryAttr", "">:$configuration
+ AttrParameter<"DictionaryAttr", "">:$configuration,
+ ArrayRefParameter<"ExecutableTargetAttr", "">:$executable_targets
);
let builders = [
AttrBuilder<(ins "StringRef":$deviceID)>,
@@ -513,9 +514,6 @@
// configuration dictionary.
bool hasConfigurationAttr(StringRef name);
- // Returns zero or more executable targets that this device supports.
- SmallVector<IREE::HAL::ExecutableTargetAttr, 4> getExecutableTargets();
-
// Returns a list of target devices that may be active for the given
// operation. This will recursively walk parent operations until one with
// the `hal.device.targets` attribute is found.
diff --git a/compiler/src/iree/compiler/Dialect/HAL/IR/test/attributes.mlir b/compiler/src/iree/compiler/Dialect/HAL/IR/test/attributes.mlir
index fd70f70..95ea673 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/IR/test/attributes.mlir
+++ b/compiler/src/iree/compiler/Dialect/HAL/IR/test/attributes.mlir
@@ -60,6 +60,20 @@
// -----
+// CHECK-LABEL: "device.targets"
+"device.targets"() {
+ // CHECK-SAME: target_0 = #hal.device.target<"a">
+ target_0 = #hal.device.target<"a">,
+ // CHECK-SAME: target_1 = #hal.device.target<"b", {config}>,
+ target_1 = #hal.device.target<"b", {config}>,
+ // CHECK-SAME: target_2 = #hal.device.target<"c", {config}, [#hal.executable.target<"llvm-cpu", "f">]>,
+ target_2 = #hal.device.target<"c", {config}, [#hal.executable.target<"llvm-cpu", "f">]>,
+ // CHECK-SAME: target_3 = #hal.device.target<"d", [#hal.executable.target<"llvm-cpu", "f">]>
+ target_3 = #hal.device.target<"d", [#hal.executable.target<"llvm-cpu", "f">]>
+} : () -> ()
+
+// -----
+
"affinity.queue"() {
// CHECK: any = #hal.affinity.queue<*>
any = #hal.affinity.queue<*>,
diff --git a/compiler/src/iree/compiler/Dialect/HAL/Target/LLVMCPU/LLVMCPUTarget.cpp b/compiler/src/iree/compiler/Dialect/HAL/Target/LLVMCPU/LLVMCPUTarget.cpp
index c4e59f5..b3dba8e 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/Target/LLVMCPU/LLVMCPUTarget.cpp
+++ b/compiler/src/iree/compiler/Dialect/HAL/Target/LLVMCPU/LLVMCPUTarget.cpp
@@ -179,12 +179,15 @@
Builder b(context);
SmallVector<NamedAttribute> configItems;
- configItems.emplace_back(b.getStringAttr("executable_targets"),
- getExecutableTargets(context, target, addlConfig));
-
auto configAttr = b.getDictionaryAttr(configItems);
+
+ // If we had multiple target environments we would generate one target attr
+ // per environment, with each setting its own environment attribute.
+ SmallVector<IREE::HAL::ExecutableTargetAttr> targetAttrs;
+ targetAttrs.push_back(getExecutableTarget(context, target, addlConfig));
+
return IREE::HAL::DeviceTargetAttr::get(
- context, b.getStringAttr(deviceID()), configAttr);
+ context, b.getStringAttr(deviceID()), configAttr, targetAttrs);
}
IREE::HAL::DeviceTargetAttr
@@ -765,15 +768,6 @@
}
private:
- ArrayAttr
- getExecutableTargets(MLIRContext *context, const LLVMTarget &target,
- const AdditionalConfigurationValues &addlConfig) const {
- SmallVector<Attribute> targetAttrs;
- // This is where we would multiversion.
- targetAttrs.push_back(getExecutableTarget(context, target, addlConfig));
- return ArrayAttr::get(context, targetAttrs);
- }
-
IREE::HAL::ExecutableTargetAttr
getExecutableTarget(MLIRContext *context, const LLVMTarget &target,
const AdditionalConfigurationValues &addlConfig) const {
diff --git a/compiler/src/iree/compiler/Dialect/HAL/Target/LLVMCPU/test/smoketest_embedded.mlir b/compiler/src/iree/compiler/Dialect/HAL/Target/LLVMCPU/test/smoketest_embedded.mlir
index 4c059a5..493a3c7 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/Target/LLVMCPU/test/smoketest_embedded.mlir
+++ b/compiler/src/iree/compiler/Dialect/HAL/Target/LLVMCPU/test/smoketest_embedded.mlir
@@ -3,11 +3,9 @@
module attributes {
hal.device.targets = [
- #hal.device.target<"llvm-cpu", {
- executable_targets = [
- #hal.executable.target<"llvm-cpu", "embedded-elf-x86_64", { native_vector_size = 16 : index }>
- ]
- }>
+ #hal.device.target<"llvm-cpu", [
+ #hal.executable.target<"llvm-cpu", "embedded-elf-x86_64", { native_vector_size = 16 : index }>
+ ]>
]
} {
diff --git a/compiler/src/iree/compiler/Dialect/HAL/Target/LLVMCPU/test/smoketest_system.mlir b/compiler/src/iree/compiler/Dialect/HAL/Target/LLVMCPU/test/smoketest_system.mlir
index a2520d2..bb5c607 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/Target/LLVMCPU/test/smoketest_system.mlir
+++ b/compiler/src/iree/compiler/Dialect/HAL/Target/LLVMCPU/test/smoketest_system.mlir
@@ -5,11 +5,9 @@
module attributes {
hal.device.targets = [
- #hal.device.target<"llvm-cpu", {
- executable_targets = [
- #hal.executable.target<"llvm-cpu", "embedded-elf-x86_64",{ native_vector_size = 16 : index } >
- ]
- }>
+ #hal.device.target<"llvm-cpu", [
+ #hal.executable.target<"llvm-cpu", "embedded-elf-x86_64",{ native_vector_size = 16 : index } >
+ ]>
]
} {
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 4b54096..e99e7bd 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/VulkanSPIRVTarget.cpp
+++ b/compiler/src/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/VulkanSPIRVTarget.cpp
@@ -128,12 +128,24 @@
Builder b(context);
SmallVector<NamedAttribute> configItems;
- configItems.emplace_back(b.getStringAttr("executable_targets"),
- getExecutableTargets(context));
-
auto configAttr = b.getDictionaryAttr(configItems);
+
+ // Select SPIR-V environments to compile for.
+ SmallVector<IREE::HAL::ExecutableTargetAttr> targetAttrs;
+ for (std::string targetTripleOrEnv : options_.targetTriplesAndEnvs) {
+ targetAttrs.push_back(getExecutableTarget(
+ context, getSPIRVTargetEnv(targetTripleOrEnv, context),
+ options_.indirectBindings));
+ }
+ // If no environment specified, populate with a minimal target.
+ if (targetAttrs.empty()) {
+ targetAttrs.push_back(getExecutableTarget(
+ context, getSPIRVTargetEnv("unknown-unknown-unknown", context),
+ options_.indirectBindings));
+ }
+
return IREE::HAL::DeviceTargetAttr::get(
- context, b.getStringAttr(deviceID()), configAttr);
+ context, b.getStringAttr(deviceID()), configAttr, targetAttrs);
}
void buildConfigurationPassPipeline(IREE::HAL::ExecutableVariantOp variantOp,
@@ -389,24 +401,6 @@
}
private:
- ArrayAttr getExecutableTargets(MLIRContext *context) const {
- SmallVector<Attribute> targetAttrs;
-
- for (std::string targetTripleOrEnv : options_.targetTriplesAndEnvs) {
- targetAttrs.push_back(getExecutableTarget(
- context, getSPIRVTargetEnv(targetTripleOrEnv, context),
- options_.indirectBindings));
- }
-
- // If no environment specified, populate with a minimal target.
- if (targetAttrs.empty()) {
- targetAttrs.push_back(getExecutableTarget(
- context, getSPIRVTargetEnv("unknown-unknown-unknown", context),
- options_.indirectBindings));
- }
- return ArrayAttr::get(context, targetAttrs);
- }
-
IREE::HAL::ExecutableTargetAttr
getExecutableTarget(MLIRContext *context, spirv::TargetEnvAttr targetEnv,
bool indirectBindings) const {
diff --git a/compiler/src/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/test/smoketest.mlir b/compiler/src/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/test/smoketest.mlir
index e51bd37..68d6542 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/test/smoketest.mlir
+++ b/compiler/src/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/test/smoketest.mlir
@@ -2,13 +2,11 @@
module attributes {
hal.device.targets = [
- #hal.device.target<"vulkan", {
- executable_targets = [
- #hal.executable.target<"vulkan-spirv", "vulkan-spirv-fb", {
- spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spirv.resource_limits<>>
- }>
- ]
- }>
+ #hal.device.target<"vulkan", [
+ #hal.executable.target<"vulkan-spirv", "vulkan-spirv-fb", {
+ spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spirv.resource_limits<>>
+ }>
+ ]>
]
} {
diff --git a/compiler/src/iree/compiler/Dialect/HAL/Transforms/test/dump_executable_benchmarks.mlir b/compiler/src/iree/compiler/Dialect/HAL/Transforms/test/dump_executable_benchmarks.mlir
index 98bcc79..b42b38f 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/Transforms/test/dump_executable_benchmarks.mlir
+++ b/compiler/src/iree/compiler/Dialect/HAL/Transforms/test/dump_executable_benchmarks.mlir
@@ -4,9 +4,9 @@
// but this is much easier to test with lit.
#executable_target_embedded_elf_x86_64 = #hal.executable.target<"llvm-cpu", "embedded-elf-x86_64">
-#device_target_cpu = #hal.device.target<"llvm-cpu", {
- executable_targets = [#executable_target_embedded_elf_x86_64]
-}>
+#device_target_cpu = #hal.device.target<"llvm-cpu", [
+ #executable_target_embedded_elf_x86_64
+]>
#pipeline_layout_0 = #hal.pipeline.layout<push_constants = 2, sets = [
#hal.descriptor_set.layout<0, bindings = [
#hal.descriptor_set.binding<0, storage_buffer>,
diff --git a/compiler/src/iree/compiler/Dialect/HAL/Transforms/test/dump_executable_sources.mlir b/compiler/src/iree/compiler/Dialect/HAL/Transforms/test/dump_executable_sources.mlir
index 6c9abfc..458e841 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/Transforms/test/dump_executable_sources.mlir
+++ b/compiler/src/iree/compiler/Dialect/HAL/Transforms/test/dump_executable_sources.mlir
@@ -4,9 +4,9 @@
// but this is much easier to test with lit.
#executable_target_embedded_elf_x86_64 = #hal.executable.target<"llvm-cpu", "embedded-elf-x86_64">
-#device_target_cpu = #hal.device.target<"llvm-cpu", {
- executable_targets = [#executable_target_embedded_elf_x86_64]
-}>
+#device_target_cpu = #hal.device.target<"llvm-cpu", [
+ #executable_target_embedded_elf_x86_64
+]>
#pipeline_layout = #hal.pipeline.layout<push_constants = 0, sets = [
#hal.descriptor_set.layout<0, bindings = [
#hal.descriptor_set.binding<0, storage_buffer>,
diff --git a/compiler/src/iree/compiler/Dialect/HAL/Transforms/test/materialize_dispatch_instrumentation.mlir b/compiler/src/iree/compiler/Dialect/HAL/Transforms/test/materialize_dispatch_instrumentation.mlir
index c51b98f..607c876 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/Transforms/test/materialize_dispatch_instrumentation.mlir
+++ b/compiler/src/iree/compiler/Dialect/HAL/Transforms/test/materialize_dispatch_instrumentation.mlir
@@ -1,12 +1,10 @@
// RUN: iree-opt --split-input-file --pass-pipeline='builtin.module(iree-hal-materialize-dispatch-instrumentation{buffer-size=64mib})' %s | FileCheck %s
module attributes {hal.device.targets = [
- #hal.device.target<"llvm-cpu", {
- executable_targets = [
- #hal.executable.target<"llvm-cpu", "embedded-elf-arm_64">,
- #hal.executable.target<"llvm-cpu", "embedded-elf-x86_64">
- ]
- }>
+ #hal.device.target<"llvm-cpu", [
+ #hal.executable.target<"llvm-cpu", "embedded-elf-arm_64">,
+ #hal.executable.target<"llvm-cpu", "embedded-elf-x86_64">
+ ]>
]} {
// Instrumentation storage buffer allocated at startup (defaults to 64MB + footer):
diff --git a/compiler/src/iree/compiler/Dialect/HAL/Transforms/test/materialize_interfaces.mlir b/compiler/src/iree/compiler/Dialect/HAL/Transforms/test/materialize_interfaces.mlir
index 87eb7a1..a688bdd 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/Transforms/test/materialize_interfaces.mlir
+++ b/compiler/src/iree/compiler/Dialect/HAL/Transforms/test/materialize_interfaces.mlir
@@ -3,12 +3,10 @@
// Tests an executable with a workgroup count region specified.
module attributes {hal.device.targets = [
- #hal.device.target<"llvm-cpu", {
- executable_targets = [
- #hal.executable.target<"llvm-cpu", "arm_64">,
- #hal.executable.target<"llvm-cpu", "x86_64">
- ]
- }>
+ #hal.device.target<"llvm-cpu", [
+ #hal.executable.target<"llvm-cpu", "arm_64">,
+ #hal.executable.target<"llvm-cpu", "x86_64">
+ ]>
]} {
// CHECK: #pipeline_layout = #hal.pipeline.layout<
// CHECK-SAME: push_constants = 1
@@ -78,12 +76,10 @@
// The default device when none is specified.
// Functions and scopes can override the target device.
hal.device.targets = [
- #hal.device.target<"cpu", {
- executable_targets = [
- #hal.executable.target<"llvm-cpu", "arm_64">,
- #hal.executable.target<"llvm-cpu", "x86_64">
- ]
- }>
+ #hal.device.target<"cpu", [
+ #hal.executable.target<"llvm-cpu", "arm_64">,
+ #hal.executable.target<"llvm-cpu", "x86_64">
+ ]>
]
} {
// CHECK: hal.executable private @ex
@@ -121,11 +117,9 @@
// CHECK-LABEL: @using_specialized
util.func public @using_specialized(%arg0: !stream.resource<transient>, %arg1: index) -> !stream.timepoint attributes {
hal.device.targets = [
- #hal.device.target<"cpu", {
- executable_targets = [
- #hal.executable.target<"llvm-cpu", "riscv_32">
- ]
- }>
+ #hal.device.target<"cpu", [
+ #hal.executable.target<"llvm-cpu", "riscv_32">
+ ]>
]
} {
%c0 = arith.constant 0 : index
@@ -153,12 +147,10 @@
// The default device when none is specified.
// Functions and scopes can override the target device.
hal.device.targets = [
- #hal.device.target<"cpu", {
- executable_targets = [
- #hal.executable.target<"llvm-cpu", "arm_64">,
- #hal.executable.target<"llvm-cpu", "x86_64">
- ]
- }>
+ #hal.device.target<"cpu", [
+ #hal.executable.target<"llvm-cpu", "arm_64">,
+ #hal.executable.target<"llvm-cpu", "x86_64">
+ ]>
]
} {
// CHECK: hal.executable private @ex
@@ -197,11 +189,9 @@
// CHECK-LABEL: @using_specialized
util.func public @using_specialized(%arg0: !stream.resource<transient>, %arg1: index) -> !stream.timepoint attributes {
hal.device.targets = [
- #hal.device.target<"cpu", {
- executable_targets = [
- #hal.executable.target<"llvm-cpu", "riscv_32">
- ]
- }>
+ #hal.device.target<"cpu", [
+ #hal.executable.target<"llvm-cpu", "riscv_32">
+ ]>
]
} {
%c0 = arith.constant 0 : index
@@ -225,12 +215,10 @@
module attributes {
hal.device.targets = [
- #hal.device.target<"cpu", {
- executable_targets = [
- #hal.executable.target<"llvm-cpu", "arm_64">,
- #hal.executable.target<"llvm-cpu", "x86_64">
- ]
- }>
+ #hal.device.target<"cpu", [
+ #hal.executable.target<"llvm-cpu", "arm_64">,
+ #hal.executable.target<"llvm-cpu", "x86_64">
+ ]>
]
} {
// CHECK: hal.executable public @ex
@@ -252,11 +240,9 @@
// CHECK-LABEL: @using_specialized
util.func public @using_specialized(%arg0: !stream.resource<transient>, %arg1: index) -> !stream.timepoint attributes {
hal.device.targets = [
- #hal.device.target<"cpu", {
- executable_targets = [
- #hal.executable.target<"llvm-cpu", "riscv_32">
- ]
- }>
+ #hal.device.target<"cpu", [
+ #hal.executable.target<"llvm-cpu", "riscv_32">
+ ]>
]
} {
%c0 = arith.constant 0 : index
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
index e1f620a..02add52 100644
--- 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
@@ -18,7 +18,7 @@
// Verify that the order of target environments matches what the user specified.
-// CHECK: executable_targets = [#[[RDNA3]], #[[ENV]], #[[VALHALL]]]
+// CHECK: [#[[RDNA3]], #[[ENV]], #[[VALHALL]]]
stream.executable public @reduce_dispatch {
stream.executable.export @reduce_dispatch workgroups(%arg0: index) -> (index, index, index) {
diff --git a/compiler/src/iree/compiler/GlobalOptimization/test/materialize_homogeneous_encodings.mlir b/compiler/src/iree/compiler/GlobalOptimization/test/materialize_homogeneous_encodings.mlir
index 9184e99..475f965 100644
--- a/compiler/src/iree/compiler/GlobalOptimization/test/materialize_homogeneous_encodings.mlir
+++ b/compiler/src/iree/compiler/GlobalOptimization/test/materialize_homogeneous_encodings.mlir
@@ -5,7 +5,7 @@
#map1 = affine_map<(d0, d1, d2) -> (d0, d2)>
#map2 = affine_map<(d0, d1, d2) -> (d2, d1)>
#map3 = affine_map<(d0, d1, d2) -> (d0, d1)>
-#device_target_llvm_cpu = #hal.device.target<"llvm-cpu", {executable_targets = [#executable_target_embedded_elf_x86_64_]}>
+#device_target_llvm_cpu = #hal.device.target<"llvm-cpu", [#executable_target_embedded_elf_x86_64_]>
module attributes {hal.device.targets = [#device_target_llvm_cpu]} {
util.func public @lhs_encoding(%arg0: tensor<?x?xf32>) -> tensor<?x?xf32> {
%cst = arith.constant 0.000000e+00 : f32
@@ -36,7 +36,7 @@
#map1 = affine_map<(d0, d1, d2) -> (d0, d2)>
#map2 = affine_map<(d0, d1, d2) -> (d2, d1)>
#map3 = affine_map<(d0, d1, d2) -> (d0, d1)>
-#device_target_vulkan = #hal.device.target<"vulkan", {executable_targets = [#executable_target_vulkan_spirv_fb]}>
+#device_target_vulkan = #hal.device.target<"vulkan", [#executable_target_vulkan_spirv_fb]>
module attributes {hal.device.targets = [#device_target_vulkan]} {
util.func public @lhs_encoding(%arg0: tensor<?x?xf32>) -> tensor<?x?xf32> {
%cst = arith.constant 0.000000e+00 : f32
@@ -69,9 +69,9 @@
#map2 = affine_map<(d0, d1, d2) -> (d2, d1)>
#map3 = affine_map<(d0, d1, d2) -> (d0, d1)>
#executable_target_embedded_elf_x86_64_ = #hal.executable.target<"llvm-cpu", "embedded-elf-x86_64", {target_triple = "x86_64-none-elf", cpu_features = "+avx512f"}>
-#device_target_llvm_cpu = #hal.device.target<"llvm-cpu", {executable_targets = [#executable_target_embedded_elf_x86_64_]}>
+#device_target_llvm_cpu = #hal.device.target<"llvm-cpu", [#executable_target_embedded_elf_x86_64_]>
#executable_target_vulkan_spirv_fb = #hal.executable.target<"vulkan", "vulkan-spirv-fb">
-#device_target_vulkan = #hal.device.target<"vulkan", {executable_targets = [#executable_target_vulkan_spirv_fb]}>
+#device_target_vulkan = #hal.device.target<"vulkan", [#executable_target_vulkan_spirv_fb]>
module attributes {hal.device.targets = [#device_target_vulkan, #device_target_llvm_cpu]} {
util.func public @lhs_encoding(%arg0: tensor<?x?xf32>) -> tensor<?x?xf32> {
%cst = arith.constant 0.000000e+00 : f32
diff --git a/samples/custom_dispatch/cpu/embedded/example_hal.mlir b/samples/custom_dispatch/cpu/embedded/example_hal.mlir
index d94d34c..e9edfd5 100644
--- a/samples/custom_dispatch/cpu/embedded/example_hal.mlir
+++ b/samples/custom_dispatch/cpu/embedded/example_hal.mlir
@@ -41,11 +41,9 @@
// These can come from compiler flags and multiple targets can be supported
// It's possible, for example, to support targeting multiple devices in the same
// compiled binary (CPU + Vulkan, etc).
-#cpu_target = #hal.device.target<"llvm-cpu", {
- executable_targets = [
- #x86_64_target
- ]
-}>
+#cpu_target = #hal.device.target<"llvm-cpu", [
+ #x86_64_target
+]>
module @example attributes {hal.device.targets = [#cpu_target]} {
diff --git a/samples/custom_dispatch/cpu/embedded/example_stream.mlir b/samples/custom_dispatch/cpu/embedded/example_stream.mlir
index 2620ca7..a8b6861 100644
--- a/samples/custom_dispatch/cpu/embedded/example_stream.mlir
+++ b/samples/custom_dispatch/cpu/embedded/example_stream.mlir
@@ -45,12 +45,10 @@
// These can come from compiler flags and multiple targets can be supported
// It's possible, for example, to support targeting multiple devices in the same
// compiled binary (CPU + Vulkan, etc).
-#cpu_target = #hal.device.target<"llvm-cpu", {
- executable_targets = [
- #arm_64_target,
- #x86_64_target
- ]
-}>
+#cpu_target = #hal.device.target<"llvm-cpu", [
+ #arm_64_target,
+ #x86_64_target
+]>
module @example attributes {hal.device.targets = [#cpu_target]} {
diff --git a/samples/custom_dispatch/cpu/embedded/example_transform.mlir b/samples/custom_dispatch/cpu/embedded/example_transform.mlir
index 0735189..709a016 100644
--- a/samples/custom_dispatch/cpu/embedded/example_transform.mlir
+++ b/samples/custom_dispatch/cpu/embedded/example_transform.mlir
@@ -26,11 +26,9 @@
// multiple targets, but this example is maintaining an implicit requirement
// that the custom kernel being spliced in is supported by the target device,
// hence we only support llvm-cpu here.
-#cpu_target = #hal.device.target<"llvm-cpu", {
- executable_targets = [
- #x86_64_target
- ]
-}>
+#cpu_target = #hal.device.target<"llvm-cpu", [
+ #x86_64_target
+]>
#map = affine_map<(d0, d1) -> (d0, d1)>
#map1 = affine_map<(d0, d1) -> (d0)>
diff --git a/samples/custom_dispatch/cpu/embedded/example_transform_spec.mlir b/samples/custom_dispatch/cpu/embedded/example_transform_spec.mlir
index b19e6d3..3726726 100644
--- a/samples/custom_dispatch/cpu/embedded/example_transform_spec.mlir
+++ b/samples/custom_dispatch/cpu/embedded/example_transform_spec.mlir
@@ -16,11 +16,9 @@
// These can come from compiler flags and multiple targets can be supported
// It's possible, for example, to support targeting multiple devices in the same
// compiled binary (CPU + Vulkan, etc).
-#cpu_target = #hal.device.target<"llvm-cpu", {
- executable_targets = [
- #x86_64_target
- ]
-}>
+#cpu_target = #hal.device.target<"llvm-cpu", [
+ #x86_64_target
+]>
module attributes {transform.with_named_sequence} {
diff --git a/samples/custom_dispatch/cpu/mlp_plugin/mlp.mlir b/samples/custom_dispatch/cpu/mlp_plugin/mlp.mlir
index 3d99d26..e7433e8 100644
--- a/samples/custom_dispatch/cpu/mlp_plugin/mlp.mlir
+++ b/samples/custom_dispatch/cpu/mlp_plugin/mlp.mlir
@@ -18,11 +18,9 @@
// multiple targets, but this example is maintaining an implicit requirement
// that the custom kernel being spliced in is supported by the target device,
// hence we only support llvm-cpu here.
-#cpu_target = #hal.device.target<"llvm-cpu", {
- executable_targets = [
- #x86_64_target
- ]
-}>
+#cpu_target = #hal.device.target<"llvm-cpu", [
+ #x86_64_target
+]>
#map = affine_map<(d0, d1) -> (d0, d1)>
module @example attributes {hal.device.targets = [#cpu_target]} {
diff --git a/samples/custom_dispatch/cpu/mlp_plugin/mlp_torch.mlir b/samples/custom_dispatch/cpu/mlp_plugin/mlp_torch.mlir
index 42a9d31..a2f8549 100644
--- a/samples/custom_dispatch/cpu/mlp_plugin/mlp_torch.mlir
+++ b/samples/custom_dispatch/cpu/mlp_plugin/mlp_torch.mlir
@@ -38,11 +38,9 @@
// multiple targets, but this example is maintaining an implicit requirement
// that the custom kernel being spliced in is supported by the target device,
// hence we only support llvm-cpu here.
-#cpu_target = #hal.device.target<"llvm-cpu", {
- executable_targets = [
- #x86_64_target
- ]
-}>
+#cpu_target = #hal.device.target<"llvm-cpu", [
+ #x86_64_target
+]>
#map = affine_map<(d0, d1) -> (d0, d1)>
module @example attributes {hal.device.targets = [#cpu_target]} {
diff --git a/samples/custom_dispatch/cpu/mlp_plugin/mlp_tosa.mlir b/samples/custom_dispatch/cpu/mlp_plugin/mlp_tosa.mlir
index 10d418e..58990ca 100644
--- a/samples/custom_dispatch/cpu/mlp_plugin/mlp_tosa.mlir
+++ b/samples/custom_dispatch/cpu/mlp_plugin/mlp_tosa.mlir
@@ -38,11 +38,9 @@
// multiple targets, but this example is maintaining an implicit requirement
// that the custom kernel being spliced in is supported by the target device,
// hence we only support llvm-cpu here.
-#cpu_target = #hal.device.target<"llvm-cpu", {
- executable_targets = [
- #x86_64_target
- ]
-}>
+#cpu_target = #hal.device.target<"llvm-cpu", [
+ #x86_64_target
+]>
module @example attributes {hal.device.targets = [#cpu_target]} {
func.func @mlp_invocation(%lhs: tensor<2x4xf32>, %rhs : tensor<4x8xf32>) -> tensor<2x8xf32> {
diff --git a/samples/custom_dispatch/cuda/kernels/example.mlir b/samples/custom_dispatch/cuda/kernels/example.mlir
index 1c64964..15a3bb43 100644
--- a/samples/custom_dispatch/cuda/kernels/example.mlir
+++ b/samples/custom_dispatch/cuda/kernels/example.mlir
@@ -24,12 +24,10 @@
// These can come from compiler flags and multiple targets can be supported
// It's possible, for example, to support targeting multiple devices in the same
// compiled binary.
-#cuda_target = #hal.device.target<"cuda", {
- executable_targets = [
- #nvptx_sm_52_target,
- #nvptx_sm_80_target
- ]
-}>
+#cuda_target = #hal.device.target<"cuda", [
+ #nvptx_sm_52_target,
+ #nvptx_sm_80_target
+]>
module @example attributes {hal.device.targets = [#cuda_target]} {
diff --git a/samples/custom_dispatch/vulkan/shaders/example.mlir b/samples/custom_dispatch/vulkan/shaders/example.mlir
index 4dd479d..d9a6b84 100644
--- a/samples/custom_dispatch/vulkan/shaders/example.mlir
+++ b/samples/custom_dispatch/vulkan/shaders/example.mlir
@@ -24,9 +24,9 @@
// These can come from compiler flags and multiple targets can be supported
// It's possible, for example, to support targeting multiple devices in the same
// compiled binary.
-#vulkan_target = #hal.device.target<"vulkan", {
- executable_targets = [#spirv_target]
-}>
+#vulkan_target = #hal.device.target<"vulkan", [
+ #spirv_target
+]>
module @example attributes {hal.device.targets = [#vulkan_target]} {
diff --git a/samples/custom_dispatch/vulkan/shaders/example_inline.mlir b/samples/custom_dispatch/vulkan/shaders/example_inline.mlir
index ea81680..9e62b8d 100644
--- a/samples/custom_dispatch/vulkan/shaders/example_inline.mlir
+++ b/samples/custom_dispatch/vulkan/shaders/example_inline.mlir
@@ -24,9 +24,7 @@
// These can come from compiler flags and multiple targets can be supported
// It's possible, for example, to support targeting multiple devices in the same
// compiled binary.
-#vulkan_target = #hal.device.target<"vulkan", {
- executable_targets = [#spirv_target]
-}>
+#vulkan_target = #hal.device.target<"vulkan", [#spirv_target]>
module @example attributes {hal.device.targets = [#vulkan_target]} {
diff --git a/samples/custom_dispatch/vulkan/shaders/example_transform.mlir b/samples/custom_dispatch/vulkan/shaders/example_transform.mlir
index d48c2f9..c562a89 100644
--- a/samples/custom_dispatch/vulkan/shaders/example_transform.mlir
+++ b/samples/custom_dispatch/vulkan/shaders/example_transform.mlir
@@ -31,9 +31,7 @@
// hence we only support vulkan here. It is possible to hand author a custom
// kernel that supports multiple targets by specifying an object per-target, but
// that requires authoring the kernel for multiple targets.
-#vulkan_target = #hal.device.target<"vulkan", {
- executable_targets = [#spirv_target]
-}>
+#vulkan_target = #hal.device.target<"vulkan", [#spirv_target]>
#map = affine_map<(d0, d1) -> (d0, d1)>
#map1 = affine_map<(d0, d1) -> (d0)>
diff --git a/samples/transform_dialect/example_module.mlir b/samples/transform_dialect/example_module.mlir
index 7b4743b..22f2927 100644
--- a/samples/transform_dialect/example_module.mlir
+++ b/samples/transform_dialect/example_module.mlir
@@ -27,7 +27,15 @@
#target_env = #spirv.target_env<#spirv.vce<v1.3, [Shader, GroupNonUniform], [SPV_KHR_storage_buffer_storage_class, SPV_KHR_variable_pointers]>, api=Vulkan, #spirv.resource_limits<max_compute_workgroup_size = [128, 128, 64], subgroup_size = 64, cooperative_matrix_properties_khr = []>>
-module attributes {hal.device.targets = [#hal.device.target<"vulkan", {executable_targets = [#hal.executable.target<"vulkan", "vulkan-spirv-fb", {spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Shader, GroupNonUniform], [SPV_KHR_storage_buffer_storage_class, SPV_KHR_variable_pointers]>, api=Vulkan, #spirv.resource_limits<max_compute_workgroup_size = [128, 128, 64], subgroup_size = 64, cooperative_matrix_properties_khr = []>>}>]}>]} {
+module attributes {
+ hal.device.targets = [
+ #hal.device.target<"vulkan", [
+ #hal.executable.target<"vulkan", "vulkan-spirv-fb", {
+ spirv.target_env = #target_env
+ }>
+ ]>
+ ]
+} {
hal.executable private @example_module_dispatch_0 {
hal.executable.variant public @vulkan_spirv_fb target(<"vulkan", "vulkan-spirv-fb", {spirv.target_env = #target_env}>) {
hal.executable.export public @example_module_dispatch_0_generic_80_f32 ordinal(0) layout(
diff --git a/tests/compiler_driver/precompile.mlir b/tests/compiler_driver/precompile.mlir
index 4e5e4fd..ecc3d0b 100644
--- a/tests/compiler_driver/precompile.mlir
+++ b/tests/compiler_driver/precompile.mlir
@@ -7,4 +7,4 @@
}
// Just check that we have the right target and executable targets.
-// CHECK: module attributes {hal.device.targets = [#hal.device.target<"vmvx", {executable_targets = [#hal.executable.target<"vmvx"
+// CHECK: module attributes {hal.device.targets = [#hal.device.target<"vmvx", [#hal.executable.target<"vmvx"
