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"