[TIL]

This blog will cover primarily what IREE does in each pipeline stage and what the resulting IR looks like, using a small vector add program.

The central design is a separation between host orchestration and device computation. IREE lowers these at different rates. The GPU kernel can already be LLVM/NVVM IR while the surrounding program still uses Stream operations. Later, the kernel becomes an opaque payload while the host program becomes VM bytecode that calls the HAL runtime.

The program and the evidence

Our input is a simple 4 elements vector add program:

func.func @simple_mul(%arg0: tensor<4xf32>,
                      %arg1: tensor<4xf32>) -> tensor<4xf32> {
  %0 = "arith.mulf"(%arg0, %arg1)
      : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32>
  return %0 : tensor<4xf32>
}

It computes C[i] = A[i] * B[i] for four elements. For inputs [1, 2, 3, 4] and [10, 20, 30, 40], the expected result is [10, 40, 90, 160].

A map of the stages

To view the result of each pipeline stage, we can use --compile-to option. Many passes may run within a stage, and an artifact normally contains several dialects at once (progressive lowering).

1. Input, ABI, and preprocessing: make the program callable

Input

It begins with func.func and tensor arithmetic. After public-function import, the function is represented using IREE’s utility dialect:

util.func public @simple_mul(%arg0: tensor<4xf32>,
                            %arg1: tensor<4xf32>) -> tensor<4xf32> {
  %0 = arith.mulf %arg0, %arg1 : tensor<4xf32>
  util.return %0 : tensor<4xf32>
}

The responsibility of input conversion is to bring supported source representations into forms the rest of the compiler understands. A PyTorch-derived program would need a supported import/export path before reaching the same general pipeline. Our hand-written MLIR already has straightforward tensor arithmetic, so this part is trivial.

ABI and Preprocessing

The next important transformation is the ABI boundary, which means separating external objects from internal ones, and define how the program interact with external environment. The public signature accepts and returns !hal.buffer_view:

%a = hal.tensor.import %arg0 "input0"
    : !hal.buffer_view -> tensor<4xf32>
%b = hal.tensor.import %arg1 "input1"
    : !hal.buffer_view -> tensor<4xf32>
%c = arith.mulf %a, %b : tensor<4xf32>
%view = hal.tensor.export %c "output0"
    : tensor<4xf32> -> !hal.buffer_view

A buffer view is the runtime-facing object that describes a buffer together with tensor metadata. Inside the function, optimization still works on tensors. Reflection retains the original signature:

sync func @simple_mul(%input0: tensor<4xf32>,
                     %input1: tensor<4xf32>) -> (%output0: tensor<4xf32>)

This separates the external object model from the internal mathematical representation. import and export mark that boundary; their presence alone does not imply a physical data copy.

Preprocessing is usually canonicalization. Since the program is small here there’s really not much to canonicalize, so the result is the same as ABI pipeline stage output. Complete preprocessing & ABI results here:

module {
  util.func public @simple_mul(%arg0: !hal.buffer_view, %arg1: !hal.buffer_view) -> !hal.buffer_view attributes {iree.abi.stub, iree.reflection = {iree.abi.declaration = "sync func @simple_mul(%input0: tensor<4xf32>, %input1: tensor<4xf32>) -> (%output0: tensor<4xf32>)"}} {
    %0 = hal.tensor.import %arg0 "input0" : !hal.buffer_view -> tensor<4xf32>
    %1 = hal.tensor.import %arg1 "input1" : !hal.buffer_view -> tensor<4xf32>
    %2 = arith.mulf %0, %1 : tensor<4xf32>
    %3 = hal.tensor.export %2 "output0" : tensor<4xf32> -> !hal.buffer_view
    util.return %3 : !hal.buffer_view
  }
}

2. Global optimization

Global optimization typically includes graph level rewrite, which means CSE, constant folding, etc. It is completely target agnostic.

In the example, the tensor-wide multiplication becomes a structured linalg.generic:

indexing_maps = [
  affine_map<(d0) -> (d0)>,
  affine_map<(d0) -> (d0)>,
  affine_map<(d0) -> (d0)>
]
iterator_types = ["parallel"]

The three maps describe how the loop index addresses input A, input B, and output C. All use the same index. The body performs scalar multiplication and yields the result:

^bb0(%a: f32, %b: f32, %out: f32):
  %product = arith.mulf %a, %b : f32
  linalg.yield %product : f32

Structured iteration makes the computation available to fusion, tiling, and distribution transformations. This artifact uses the first input tensor as the destination-style outs operand. Because the scalar body ignores %out, it does not depend on the old output contents. Tensor semantics at this stage do not establish an in-place update of the caller’s input buffer; later aliasing and bufferization decisions must preserve program semantics.

The surrounding iree_tensor_ext.compute_barrier.start/end operations constrain compiler transformations, such as reshape propagation across a boundary. They are only compiler level boundary, so not materialized and is possible to not materialize.

Global optimization works while the program still exposes tensor structure. For a larger model this is where transformations can operate across a broader computation graph, before dispatch boundaries limit which operations can be combined.

The result of global optimization:

module {
  util.func public @simple_mul(%arg0: !hal.buffer_view, %arg1: !hal.buffer_view) -> !hal.buffer_view attributes {iree.abi.stub, iree.reflection = {iree.abi.declaration = "sync func @simple_mul(%input0: tensor<4xf32>, %input1: tensor<4xf32>) -> (%output0: tensor<4xf32>)"}} {
    %0 = hal.tensor.import %arg0 "input0" : !hal.buffer_view -> tensor<4xf32>
    %1 = iree_tensor_ext.compute_barrier.start %0 : tensor<4xf32> -> tensor<4xf32>
    %2 = hal.tensor.import %arg1 "input1" : !hal.buffer_view -> tensor<4xf32>
    %3 = iree_tensor_ext.compute_barrier.start %2 : tensor<4xf32> -> tensor<4xf32>
    %4 = linalg.generic {indexing_maps = [affine_map<(d0) -> (d0)>, affine_map<(d0) -> (d0)>, affine_map<(d0) -> (d0)>], iterator_types = ["parallel"]} ins(%1, %3 : tensor<4xf32>, tensor<4xf32>) outs(%1 : tensor<4xf32>) {
    ^bb0(%in: f32, %in_0: f32, %out: f32):
      %7 = arith.mulf %in, %in_0 : f32
      linalg.yield %7 : f32
    } -> tensor<4xf32>
    %5 = iree_tensor_ext.compute_barrier.end %4 : tensor<4xf32> -> tensor<4xf32>
    %6 = hal.tensor.export %5 "output0" : tensor<4xf32> -> !hal.buffer_view
    util.return %6 : !hal.buffer_view
  }
}

3. Flow

Flow is essentially doing “dispatch creation”, which decides which operations belong in a unit of executable work. Before flow, all we have is a compute graph: it does not differentiate what computation needs to grouped to a single device for SPMD. A dispatch region, represented as a “flow.executable”, is the atom for scheduling. For example, elementwise operations are natural for SPMD, so they will be submitted as a single dispatch. Reduction is not, since it changes the boundary (grid and block shapes will be changed).

After flow pipeline, the outer function now calls:

%c = flow.dispatch
    @simple_mul_dispatch_0::@simple_mul_dispatch_0_elementwise_4_f32(%a, %b)
    : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32>

The computation lives in a flow.executable, with an exported entry point and a nested function. That function receives two read-only dispatch tensors and one write-only output dispatch tensor. It loads the input slices, performs the Linalg computation, and stores the result.

This establishes a useful separation: the outer program says which computation to invoke, while the executable contains how to compute its result. A dispatch boundary is a scheduling and compilation boundary.

The export still contains:

%x, %y, %z = iree_tensor_ext.dispatch.workgroup_count_from_slice()

This is a placeholder for workgroup-count calculation. A workgroup is a unit of distributed execution; on CUDA it corresponds to a thread block. The number of groups and the number of threads per group are different quantities, and neither follows directly from the four-element tensor shape at this point.

Full result after flow pipeline stage:

module {
  flow.executable private @simple_mul_dispatch_0 {
    flow.executable.export public @simple_mul_dispatch_0_elementwise_4_f32 workgroups() -> (index, index, index) {
      %x, %y, %z = iree_tensor_ext.dispatch.workgroup_count_from_slice()
      flow.return %x, %y, %z : index, index, index
    }
    builtin.module {
      func.func @simple_mul_dispatch_0_elementwise_4_f32(%arg0: !iree_tensor_ext.dispatch.tensor<readonly:tensor<4xf32>>, %arg1: !iree_tensor_ext.dispatch.tensor<readonly:tensor<4xf32>>, %arg2: !iree_tensor_ext.dispatch.tensor<writeonly:tensor<4xf32>>) {
        %0 = iree_tensor_ext.dispatch.tensor.load %arg0, offsets = [0], sizes = [4], strides = [1] : !iree_tensor_ext.dispatch.tensor<readonly:tensor<4xf32>> -> tensor<4xf32>
        %1 = iree_tensor_ext.dispatch.tensor.load %arg1, offsets = [0], sizes = [4], strides = [1] : !iree_tensor_ext.dispatch.tensor<readonly:tensor<4xf32>> -> tensor<4xf32>
        %2 = tensor.empty() : tensor<4xf32>
        %3 = linalg.generic {indexing_maps = [affine_map<(d0) -> (d0)>, affine_map<(d0) -> (d0)>, affine_map<(d0) -> (d0)>], iterator_types = ["parallel"]} ins(%0, %1 : tensor<4xf32>, tensor<4xf32>) outs(%2 : tensor<4xf32>) {
        ^bb0(%in: f32, %in_0: f32, %out: f32):
          %4 = arith.mulf %in, %in_0 : f32
          linalg.yield %4 : f32
        } -> tensor<4xf32>
        iree_tensor_ext.dispatch.tensor.store %3, %arg2, offsets = [0], sizes = [4], strides = [1] : tensor<4xf32> -> !iree_tensor_ext.dispatch.tensor<writeonly:tensor<4xf32>>
        return
      }
    }
  }
  util.func public @simple_mul(%arg0: !hal.buffer_view, %arg1: !hal.buffer_view) -> !hal.buffer_view attributes {iree.abi.stub, iree.reflection = {iree.abi.declaration = "sync func @simple_mul(%input0: tensor<4xf32>, %input1: tensor<4xf32>) -> (%output0: tensor<4xf32>)"}} {
    %0 = hal.tensor.import %arg0 "input0" : !hal.buffer_view -> tensor<4xf32>
    %1 = hal.tensor.import %arg1 "input1" : !hal.buffer_view -> tensor<4xf32>
    %2 = flow.dispatch @simple_mul_dispatch_0::@simple_mul_dispatch_0_elementwise_4_f32(%0, %1) : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32>
    %3 = hal.tensor.export %2 "output0" : tensor<4xf32> -> !hal.buffer_view
    util.return %3 : !hal.buffer_view
  }
}

4. Stream

Stream pipeline makes resource sizes and scheduling visible. For example, when dispatching a workload, we need to answer the following question:

• Where is the input of the dispatch? Is it ready?
• Where can we execute this dispatch?
• When can host read back the result? To answer those questions, we need to track resource lifetime with allocation and dependencies.

The stream dialect rewrites the external resource access into stream.tensor.import, with explicity buffer size:

%a = stream.tensor.import %arg0
    : !hal.buffer_view -> tensor<4xf32>
      in !stream.resource<external>{%c16}

The external resource category describes a lifetime that crosses program boundaries (for example, allocated on host and copied to device, or from another kernel).

The output is explicitly allocated:

%result, %ready = stream.resource.alloca uninitialized
    : !stream.resource<external>{%c16} => !stream.timepoint

Allocation returns both a resource and a timepoint describing when it can be used. This alloca is a Stream resource operation, not a promise of CPU stack allocation.

Execution then waits for that resource and declares its accesses:

stream.cmd.dispatch @simple_mul_dispatch_0::@... {
  ro %a[0 for 16] : ...,
  ro %b[0 for 16] : ...,
  wo %result[0 for 16] : ...
}

It annotates the input parameter with two 16-byte reads and one 16-byte write. The surrounding stream.cmd.execute returns a completion timepoint; stream.timepoint.await associates subsequent use of the result with completion of that execution.

Stream separates data dependencies from a particular queue or synchronization API. An earlier async-form Stream program may contain stream.async.transfer and stream.async.dispatch, with dependencies expressed through resource values. Scheduling and allocation lower that representation toward command regions and timepoints. stream.async.transfer, when present, represents a transition between resource states or affinities. It may require copying between memory spaces, but optimization can rewrite or remove transfers.

Inside the executable, !stream.binding parameters and stream.binding.subspan describe how to obtain tensor views from bound resources. The Linalg computation itself is still high level. IREE can reason about resource scheduling before committing to the final implementation of each kernel.

Full result of Stream pipeline:

module {
  stream.executable private @simple_mul_dispatch_0 {
    stream.executable.export public @simple_mul_dispatch_0_elementwise_4_f32 workgroups() -> (index, index, index) {
      %x, %y, %z = iree_tensor_ext.dispatch.workgroup_count_from_slice()
      stream.return %x, %y, %z : index, index, index
    }
    builtin.module {
      func.func @simple_mul_dispatch_0_elementwise_4_f32(%arg0: !stream.binding {stream.alignment = 64 : index}, %arg1: !stream.binding {stream.alignment = 64 : index}, %arg2: !stream.binding {stream.alignment = 64 : index}) {
        %c0 = arith.constant 0 : index
        %0 = stream.binding.subspan %arg0[%c0] : !stream.binding -> !iree_tensor_ext.dispatch.tensor<readonly:tensor<4xf32>>
        %1 = stream.binding.subspan %arg1[%c0] : !stream.binding -> !iree_tensor_ext.dispatch.tensor<readonly:tensor<4xf32>>
        %2 = stream.binding.subspan %arg2[%c0] : !stream.binding -> !iree_tensor_ext.dispatch.tensor<writeonly:tensor<4xf32>>
        %3 = iree_tensor_ext.dispatch.tensor.load %0, offsets = [0], sizes = [4], strides = [1] : !iree_tensor_ext.dispatch.tensor<readonly:tensor<4xf32>> -> tensor<4xf32>
        %4 = iree_tensor_ext.dispatch.tensor.load %1, offsets = [0], sizes = [4], strides = [1] : !iree_tensor_ext.dispatch.tensor<readonly:tensor<4xf32>> -> tensor<4xf32>
        %5 = tensor.empty() : tensor<4xf32>
        %6 = linalg.generic {indexing_maps = [affine_map<(d0) -> (d0)>, affine_map<(d0) -> (d0)>, affine_map<(d0) -> (d0)>], iterator_types = ["parallel"]} ins(%3, %4 : tensor<4xf32>, tensor<4xf32>) outs(%5 : tensor<4xf32>) {
        ^bb0(%in: f32, %in_0: f32, %out: f32):
          %7 = arith.mulf %in, %in_0 : f32
          linalg.yield %7 : f32
        } -> tensor<4xf32>
        iree_tensor_ext.dispatch.tensor.store %6, %2, offsets = [0], sizes = [4], strides = [1] : tensor<4xf32> -> !iree_tensor_ext.dispatch.tensor<writeonly:tensor<4xf32>>
        return
      }
    }
  }
  util.func public @simple_mul(%arg0: !hal.buffer_view, %arg1: !hal.buffer_view) -> !hal.buffer_view attributes {iree.abi.stub, iree.reflection = {iree.abi.declaration = "sync func @simple_mul(%input0: tensor<4xf32>, %input1: tensor<4xf32>) -> (%output0: tensor<4xf32>)"}} {
    %c0 = arith.constant 0 : index
    %c16 = arith.constant 16 : index
    %c4 = arith.constant 4 : index
    %element_type_f32 = hal.element_type<f32> : i32
    %dense_row_major = hal.encoding_type<dense_row_major> : i32
    hal.buffer_view.assert<%arg0 : !hal.buffer_view> message("input0") shape([%c4]) type(%element_type_f32) encoding(%dense_row_major)
    %0 = stream.tensor.import %arg0 : !hal.buffer_view -> tensor<4xf32> in !stream.resource<external>{%c16}
    hal.buffer_view.assert<%arg1 : !hal.buffer_view> message("input1") shape([%c4]) type(%element_type_f32) encoding(%dense_row_major)
    %1 = stream.tensor.import %arg1 : !hal.buffer_view -> tensor<4xf32> in !stream.resource<external>{%c16}
    %result, %result_timepoint = stream.resource.alloca uninitialized : !stream.resource<external>{%c16} => !stream.timepoint
    %2 = stream.cmd.execute await(%result_timepoint) => with(%0 as %arg2: !stream.resource<external>{%c16}, %1 as %arg3: !stream.resource<external>{%c16}, %result as %arg4: !stream.resource<external>{%c16}) {
      stream.cmd.dispatch @simple_mul_dispatch_0::@simple_mul_dispatch_0_elementwise_4_f32 {
        ro %arg2[%c0 for %c16] : !stream.resource<external>{%c16},
        ro %arg3[%c0 for %c16] : !stream.resource<external>{%c16},
        wo %arg4[%c0 for %c16] : !stream.resource<external>{%c16}
      }
    } => !stream.timepoint
    %3 = stream.timepoint.await %2 => %result : !stream.resource<external>{%c16}
    %4 = stream.tensor.export %3 : tensor<4xf32> in !stream.resource<external>{%c16} -> !hal.buffer_view
    util.return %4 : !hal.buffer_view
  }
}

5. Executable sources

After this pipeline, the program has a default device affinity and a device declaration targeting CUDA. Operations such as allocation and execution are associated with that device.

The executable becomes:

hal.executable @simple_mul_dispatch_0
  └─ hal.executable.variant @cuda_nvptx_fb
       ├─ hal.executable.export
       └─ builtin.module containing the kernel

A variant is an implementation for an executable target. The structure allows multiple implementations, even though this example has only one. The export establishes a pipeline layout with three storage-buffer bindings:

The kernel obtains data with hal.interface.binding.subspan binding(0), and similarly for bindings 1 and 2. Device metadata records architectural capabilities: subgroup size choices, maximum block dimensions, and available workgroup memory. These are constraints for code generation.

It does not touch kernel internals.

module attributes {stream.affinity.default = #hal.device.affinity<@__device_0>} {
  util.global private @__device_0 = #hal.device.target<"cuda", [#hal.executable.target<"cuda", "cuda-nvptx-fb", {iree_codegen.target_info = #iree_gpu.target<arch = "sm_60", features = "+ptx76", wgp = <compute = fp64|fp32|fp16|int64|int32|int16|int8, storage = b64|b32|b16|b8, subgroup = shuffle|arithmetic, subgroup_size_choices = [32], max_workgroup_sizes = [1024, 1024, 1024], max_thread_count_per_workgroup = 1024, max_workgroup_memory_bytes = 49152, max_workgroup_counts = [2147483647, 65535, 65535]>>}>]> : !hal.device
  hal.executable private @simple_mul_dispatch_0 {
    hal.executable.variant public @cuda_nvptx_fb target(<"cuda", "cuda-nvptx-fb", {iree_codegen.target_info = #iree_gpu.target<arch = "sm_60", features = "+ptx76", wgp = <compute = fp64|fp32|fp16|int64|int32|int16|int8, storage = b64|b32|b16|b8, subgroup = shuffle|arithmetic, subgroup_size_choices = [32], max_workgroup_sizes = [1024, 1024, 1024], max_thread_count_per_workgroup = 1024, max_workgroup_memory_bytes = 49152, max_workgroup_counts = [2147483647, 65535, 65535]>>}>) {
      hal.executable.export public @simple_mul_dispatch_0_elementwise_4_f32 ordinal(0) layout(#hal.pipeline.layout<bindings = [#hal.pipeline.binding<storage_buffer, "ReadOnly|Indirect">, #hal.pipeline.binding<storage_buffer, "ReadOnly|Indirect">, #hal.pipeline.binding<storage_buffer, Indirect>], flags = Indirect>) count(%arg0: !hal.device) -> (index, index, index) {
        %x, %y, %z = iree_tensor_ext.dispatch.workgroup_count_from_slice()
        hal.return %x, %y, %z : index, index, index
      }
      builtin.module {
        func.func @simple_mul_dispatch_0_elementwise_4_f32() {
          %c0 = arith.constant 0 : index
          %0 = hal.interface.binding.subspan layout(<bindings = [#hal.pipeline.binding<storage_buffer, "ReadOnly|Indirect">, #hal.pipeline.binding<storage_buffer, "ReadOnly|Indirect">, #hal.pipeline.binding<storage_buffer, Indirect>], flags = Indirect>) binding(0) alignment(64) offset(%c0) flags("ReadOnly|Indirect") : !iree_tensor_ext.dispatch.tensor<readonly:tensor<4xf32>>
          %1 = hal.interface.binding.subspan layout(<bindings = [#hal.pipeline.binding<storage_buffer, "ReadOnly|Indirect">, #hal.pipeline.binding<storage_buffer, "ReadOnly|Indirect">, #hal.pipeline.binding<storage_buffer, Indirect>], flags = Indirect>) binding(1) alignment(64) offset(%c0) flags("ReadOnly|Indirect") : !iree_tensor_ext.dispatch.tensor<readonly:tensor<4xf32>>
          %2 = hal.interface.binding.subspan layout(<bindings = [#hal.pipeline.binding<storage_buffer, "ReadOnly|Indirect">, #hal.pipeline.binding<storage_buffer, "ReadOnly|Indirect">, #hal.pipeline.binding<storage_buffer, Indirect>], flags = Indirect>) binding(2) alignment(64) offset(%c0) flags(Indirect) : !iree_tensor_ext.dispatch.tensor<writeonly:tensor<4xf32>>
          %3 = iree_tensor_ext.dispatch.tensor.load %0, offsets = [0], sizes = [4], strides = [1] : !iree_tensor_ext.dispatch.tensor<readonly:tensor<4xf32>> -> tensor<4xf32>
          %4 = iree_tensor_ext.dispatch.tensor.load %1, offsets = [0], sizes = [4], strides = [1] : !iree_tensor_ext.dispatch.tensor<readonly:tensor<4xf32>> -> tensor<4xf32>
          %5 = tensor.empty() : tensor<4xf32>
          %6 = linalg.generic {indexing_maps = [affine_map<(d0) -> (d0)>, affine_map<(d0) -> (d0)>, affine_map<(d0) -> (d0)>], iterator_types = ["parallel"]} ins(%3, %4 : tensor<4xf32>, tensor<4xf32>) outs(%5 : tensor<4xf32>) {
          ^bb0(%in: f32, %in_0: f32, %out: f32):
            %7 = arith.mulf %in, %in_0 : f32
            linalg.yield %7 : f32
          } -> tensor<4xf32>
          iree_tensor_ext.dispatch.tensor.store %6, %2, offsets = [0], sizes = [4], strides = [1] : tensor<4xf32> -> !iree_tensor_ext.dispatch.tensor<writeonly:tensor<4xf32>>
          return
        }
      }
    }
  }
  util.func public @simple_mul(%arg0: !hal.buffer_view, %arg1: !hal.buffer_view) -> !hal.buffer_view attributes {iree.abi.stub, iree.reflection = {iree.abi.declaration = "sync func @simple_mul(%input0: tensor<4xf32>, %input1: tensor<4xf32>) -> (%output0: tensor<4xf32>)"}} {
    %c0 = arith.constant 0 : index
    %c16 = arith.constant 16 : index
    %c4 = arith.constant 4 : index
    %element_type_f32 = hal.element_type<f32> : i32
    %dense_row_major = hal.encoding_type<dense_row_major> : i32
    hal.buffer_view.assert<%arg0 : !hal.buffer_view> message("input0") shape([%c4]) type(%element_type_f32) encoding(%dense_row_major)
    %0 = stream.tensor.import on(#hal.device.affinity<@__device_0>) %arg0 : !hal.buffer_view -> tensor<4xf32> in !stream.resource<external>{%c16}
    hal.buffer_view.assert<%arg1 : !hal.buffer_view> message("input1") shape([%c4]) type(%element_type_f32) encoding(%dense_row_major)
    %1 = stream.tensor.import on(#hal.device.affinity<@__device_0>) %arg1 : !hal.buffer_view -> tensor<4xf32> in !stream.resource<external>{%c16}
    %result, %result_timepoint = stream.resource.alloca uninitialized on(#hal.device.affinity<@__device_0>) : !stream.resource<external>{%c16} => !stream.timepoint
    %2 = stream.cmd.execute on(#hal.device.affinity<@__device_0>) await(%result_timepoint) => with(%0 as %arg2: !stream.resource<external>{%c16}, %1 as %arg3: !stream.resource<external>{%c16}, %result as %arg4: !stream.resource<external>{%c16}) {
      stream.cmd.dispatch @simple_mul_dispatch_0::@cuda_nvptx_fb::@simple_mul_dispatch_0_elementwise_4_f32 {
        ro %arg2[%c0 for %c16] : !stream.resource<external>{%c16},
        ro %arg3[%c0 for %c16] : !stream.resource<external>{%c16},
        wo %arg4[%c0 for %c16] : !stream.resource<external>{%c16}
      }
    } => !stream.timepoint
    %3 = stream.timepoint.await %2 => %result : !stream.resource<external>{%c16}
    %4 = stream.tensor.export on(#hal.device.affinity<@__device_0>) %3 : tensor<4xf32> in !stream.resource<external>{%c16} -> !hal.buffer_view
    util.return %4 : !hal.buffer_view
  }
}

6. Executable configurations: choose a lowering strategy

It adds the most compact explanation of the intended GPU mapping:

translation_info = #iree_codegen.translation_info<
  pipeline = #iree_gpu.pipeline<TileAndFuse>
  workgroup_size = [32, 1, 1]
  subgroup_size = 32
>

The Linalg operation receives:

lowering_config = #iree_gpu.lowering_config<{
  thread = [1], workgroup = [32]
}>

These attributes answer different questions:

• workgroup_size specifies a three-dimensional thread organization: 32 threads per block.
• workgroup = [32] tiles the one-dimensional iteration space into workgroup-sized chunks of 32 elements.
• thread = [1] gives a thread-level tile of one element.
• subgroup_size = 32 selects a CUDA warp-sized subgroup.

The four-element problem therefore fits in one block, with only four threads needing to perform useful multiplication. At this boundary, that is the mapping implied by the configuration; the next artifact shows the generated control flow.

TileAndFuse selects a code generation pipeline. Its name does not imply that multiple operations were fused in this one-operation example. Keeping strategy selection separate from executing the transformations makes configurations inspectable and provides a place for user or tuning choices.

The buffer interface is also normalized. Binding subspans now produce memref<4xf32, ...> values, with iree_codegen.load_from_buffer and store_to_buffer bridging to the remaining tensor computation. Full bufferization has not happened merely because memrefs appear at the boundary.

module attributes {stream.affinity.default = #hal.device.affinity<@__device_0>} {
  util.global private @__device_0 = #hal.device.target<"cuda", [#hal.executable.target<"cuda", "cuda-nvptx-fb", {iree_codegen.target_info = #iree_gpu.target<arch = "sm_60", features = "+ptx76", wgp = <compute = fp64|fp32|fp16|int64|int32|int16|int8, storage = b64|b32|b16|b8, subgroup = shuffle|arithmetic, subgroup_size_choices = [32], max_workgroup_sizes = [1024, 1024, 1024], max_thread_count_per_workgroup = 1024, max_workgroup_memory_bytes = 49152, max_workgroup_counts = [2147483647, 65535, 65535]>>}>]> : !hal.device
  hal.executable private @simple_mul_dispatch_0 {
    hal.executable.variant public @cuda_nvptx_fb target(<"cuda", "cuda-nvptx-fb", {iree_codegen.target_info = #iree_gpu.target<arch = "sm_60", features = "+ptx76", wgp = <compute = fp64|fp32|fp16|int64|int32|int16|int8, storage = b64|b32|b16|b8, subgroup = shuffle|arithmetic, subgroup_size_choices = [32], max_workgroup_sizes = [1024, 1024, 1024], max_thread_count_per_workgroup = 1024, max_workgroup_memory_bytes = 49152, max_workgroup_counts = [2147483647, 65535, 65535]>>}>) {
      hal.executable.export public @simple_mul_dispatch_0_elementwise_4_f32 ordinal(0) layout(#hal.pipeline.layout<bindings = [#hal.pipeline.binding<storage_buffer, "ReadOnly|Indirect">, #hal.pipeline.binding<storage_buffer, "ReadOnly|Indirect">, #hal.pipeline.binding<storage_buffer, Indirect>], flags = Indirect>) count(%arg0: !hal.device) -> (index, index, index) {
        %x, %y, %z = iree_tensor_ext.dispatch.workgroup_count_from_slice()
        hal.return %x, %y, %z : index, index, index
      }
      builtin.module {
        func.func @simple_mul_dispatch_0_elementwise_4_f32() attributes {translation_info = #iree_codegen.translation_info<pipeline = #iree_gpu.pipeline<TileAndFuse> workgroup_size = [32, 1, 1] subgroup_size = 32>} {
          %c0 = arith.constant 0 : index
          %0 = hal.interface.binding.subspan layout(<bindings = [#hal.pipeline.binding<storage_buffer, "ReadOnly|Indirect">, #hal.pipeline.binding<storage_buffer, "ReadOnly|Indirect">, #hal.pipeline.binding<storage_buffer, Indirect>], flags = Indirect>) binding(0) alignment(64) offset(%c0) flags("ReadOnly|Indirect") : memref<4xf32, #hal.descriptor_type<storage_buffer>>
          %1 = hal.interface.binding.subspan layout(<bindings = [#hal.pipeline.binding<storage_buffer, "ReadOnly|Indirect">, #hal.pipeline.binding<storage_buffer, "ReadOnly|Indirect">, #hal.pipeline.binding<storage_buffer, Indirect>], flags = Indirect>) binding(1) alignment(64) offset(%c0) flags("ReadOnly|Indirect") : memref<4xf32, #hal.descriptor_type<storage_buffer>>
          %2 = hal.interface.binding.subspan layout(<bindings = [#hal.pipeline.binding<storage_buffer, "ReadOnly|Indirect">, #hal.pipeline.binding<storage_buffer, "ReadOnly|Indirect">, #hal.pipeline.binding<storage_buffer, Indirect>], flags = Indirect>) binding(2) alignment(64) offset(%c0) flags(Indirect) : memref<4xf32, #hal.descriptor_type<storage_buffer>>
          %3 = iree_codegen.load_from_buffer %0 : memref<4xf32, #hal.descriptor_type<storage_buffer>> -> tensor<4xf32>
          %4 = iree_codegen.load_from_buffer %1 : memref<4xf32, #hal.descriptor_type<storage_buffer>> -> tensor<4xf32>
          %5 = tensor.empty() : tensor<4xf32>
          %6 = linalg.generic {indexing_maps = [affine_map<(d0) -> (d0)>, affine_map<(d0) -> (d0)>, affine_map<(d0) -> (d0)>], iterator_types = ["parallel"]} ins(%3, %4 : tensor<4xf32>, tensor<4xf32>) outs(%5 : tensor<4xf32>) attrs =  {lowering_config = #iree_gpu.lowering_config<{thread = [1], workgroup = [32]}>} {
          ^bb0(%in: f32, %in_0: f32, %out: f32):
            %7 = arith.mulf %in, %in_0 : f32
            linalg.yield %7 : f32
          } -> tensor<4xf32>
          iree_codegen.store_to_buffer %6, %2 : tensor<4xf32> into memref<4xf32, #hal.descriptor_type<storage_buffer>>
          return
        }
        iree_codegen.dispatch_config @simple_mul_dispatch_0_elementwise_4_f32 {
        ^bb0(%arg0: !hal.device):
          %x, %y, %z = iree_tensor_ext.dispatch.workgroup_count_from_slice()
          iree_codegen.yield %x, %y, %z : index, index, index
        }
      }
    }
  }
  util.func public @simple_mul(%arg0: !hal.buffer_view, %arg1: !hal.buffer_view) -> !hal.buffer_view attributes {iree.abi.stub, iree.reflection = {iree.abi.declaration = "sync func @simple_mul(%input0: tensor<4xf32>, %input1: tensor<4xf32>) -> (%output0: tensor<4xf32>)"}} {
    %c0 = arith.constant 0 : index
    %c16 = arith.constant 16 : index
    %c4 = arith.constant 4 : index
    %element_type_f32 = hal.element_type<f32> : i32
    %dense_row_major = hal.encoding_type<dense_row_major> : i32
    hal.buffer_view.assert<%arg0 : !hal.buffer_view> message("input0") shape([%c4]) type(%element_type_f32) encoding(%dense_row_major)
    %0 = stream.tensor.import on(#hal.device.affinity<@__device_0>) %arg0 : !hal.buffer_view -> tensor<4xf32> in !stream.resource<external>{%c16}
    hal.buffer_view.assert<%arg1 : !hal.buffer_view> message("input1") shape([%c4]) type(%element_type_f32) encoding(%dense_row_major)
    %1 = stream.tensor.import on(#hal.device.affinity<@__device_0>) %arg1 : !hal.buffer_view -> tensor<4xf32> in !stream.resource<external>{%c16}
    %result, %result_timepoint = stream.resource.alloca uninitialized on(#hal.device.affinity<@__device_0>) : !stream.resource<external>{%c16} => !stream.timepoint
    %2 = stream.cmd.execute on(#hal.device.affinity<@__device_0>) await(%result_timepoint) => with(%0 as %arg2: !stream.resource<external>{%c16}, %1 as %arg3: !stream.resource<external>{%c16}, %result as %arg4: !stream.resource<external>{%c16}) {
      stream.cmd.dispatch @simple_mul_dispatch_0::@cuda_nvptx_fb::@simple_mul_dispatch_0_elementwise_4_f32 {
        ro %arg2[%c0 for %c16] : !stream.resource<external>{%c16},
        ro %arg3[%c0 for %c16] : !stream.resource<external>{%c16},
        wo %arg4[%c0 for %c16] : !stream.resource<external>{%c16}
      }
    } => !stream.timepoint
    %3 = stream.timepoint.await %2 => %result : !stream.resource<external>{%c16}
    %4 = stream.tensor.export on(#hal.device.affinity<@__device_0>) %3 : tensor<4xf32> in !stream.resource<external>{%c16} -> !hal.buffer_view
    util.return %4 : !hal.buffer_view
  }
}

7. Executable targets: make the GPU implementation concrete

In executable target pipeline, tensor operations, Linalg, and the configuration-only function body have been replaced by LLVM and NVVM operations. The entry now takes three !llvm.ptr<1> parameters. For this NVVM target, address space 1 is global memory. The export’s workgroup-count region returns (1, 1, 1), and its workgroup size is (32, 1, 1).

The kernel reads its thread ID:

%tid = nvvm.read.ptx.sreg.tid.x range <i32, 0, 32> : i32

The resulting control flow is equivalent to:

for (int i = threadIdx.x; i < 4; i += 32) {
  C[i] = A[i] * B[i];
}

Threads 0–3 each process one element. Threads 4–31 fail the initial bounds check and exit. No thread accesses an out-of-bounds element.

llvm.getelementptr computes element addresses; llvm.load, llvm.fmul, and llvm.store implement the arithmetic. The intermediate vector<1xf32> is a one-element vector representation, not a four-element vector operation per thread.

Pointer attributes such as readonly, noalias, and alignment information support subsequent optimization. There is no shared-memory cooperation or thread barrier because these multiplications are independent.

This boundary contains target-oriented MLIR, not PTX text yet. Meanwhile, the outer function still contains Stream operations. This is the clearest demonstration that kernel lowering and orchestration lowering proceed separately.

module attributes {stream.affinity.default = #hal.device.affinity<@__device_0>} {
  util.global private @__device_0 = #hal.device.target<"cuda", [#hal.executable.target<"cuda", "cuda-nvptx-fb", {iree_codegen.target_info = #iree_gpu.target<arch = "sm_60", features = "+ptx76", wgp = <compute = fp64|fp32|fp16|int64|int32|int16|int8, storage = b64|b32|b16|b8, subgroup = shuffle|arithmetic, subgroup_size_choices = [32], max_workgroup_sizes = [1024, 1024, 1024], max_thread_count_per_workgroup = 1024, max_workgroup_memory_bytes = 49152, max_workgroup_counts = [2147483647, 65535, 65535]>>}>]> : !hal.device
  hal.executable private @simple_mul_dispatch_0 {
    hal.executable.variant public @cuda_nvptx_fb target(<"cuda", "cuda-nvptx-fb", {iree_codegen.target_info = #iree_gpu.target<arch = "sm_60", features = "+ptx76", wgp = <compute = fp64|fp32|fp16|int64|int32|int16|int8, storage = b64|b32|b16|b8, subgroup = shuffle|arithmetic, subgroup_size_choices = [32], max_workgroup_sizes = [1024, 1024, 1024], max_thread_count_per_workgroup = 1024, max_workgroup_memory_bytes = 49152, max_workgroup_counts = [2147483647, 65535, 65535]>>}>) {
      hal.executable.export public @simple_mul_dispatch_0_elementwise_4_f32 ordinal(0) layout(#hal.pipeline.layout<bindings = [#hal.pipeline.binding<storage_buffer, "ReadOnly|Indirect">, #hal.pipeline.binding<storage_buffer, "ReadOnly|Indirect">, #hal.pipeline.binding<storage_buffer, Indirect>], flags = Indirect>) count(%arg0: !hal.device) -> (index, index, index) {
        %c1 = arith.constant 1 : index
        hal.return %c1, %c1, %c1 : index, index, index
      } attributes {subgroup_size = 32 : index, workgroup_size = [32 : index, 1 : index, 1 : index]}
      builtin.module {
        llvm.func @simple_mul_dispatch_0_elementwise_4_f32(%arg0: !llvm.ptr<1> {llvm.align = 16 : i32, llvm.noalias, llvm.nonnull, llvm.noundef, llvm.readonly}, %arg1: !llvm.ptr<1> {llvm.align = 16 : i32, llvm.noalias, llvm.nonnull, llvm.noundef, llvm.readonly}, %arg2: !llvm.ptr<1> {llvm.align = 16 : i32, llvm.noalias, llvm.nonnull, llvm.noundef}) attributes {gpu.known_block_size = array<i32: 32, 1, 1>} {
          %0 = llvm.mlir.constant(64 : i64) : i64
          %1 = llvm.mlir.constant(true) : i1
          %2 = llvm.mlir.constant(32 : i32) : i32
          %3 = llvm.mlir.constant(4 : i32) : i32
          %4 = nvvm.read.ptx.sreg.tid.x range <i32, 0, 32> : i32
          %5 = llvm.sext %4 : i32 to i64
          llvm.intr.assume %1 ["align"(%arg0, %0 : !llvm.ptr<1>, i64)] : i1
          llvm.intr.assume %1 ["align"(%arg1, %0 : !llvm.ptr<1>, i64)] : i1
          llvm.intr.assume %1 ["align"(%arg2, %0 : !llvm.ptr<1>, i64)] : i1
          %6 = llvm.trunc %5 : i64 to i32
          llvm.br ^bb1(%6 : i32)
        ^bb1(%7: i32):  // 2 preds: ^bb0, ^bb2
          %8 = llvm.icmp "slt" %7, %3 : i32
          llvm.cond_br %8, ^bb2, ^bb3
        ^bb2:  // pred: ^bb1
          %9 = llvm.zext %7 : i32 to i64
          %10 = llvm.getelementptr %arg0[%9] : (!llvm.ptr<1>, i64) -> !llvm.ptr<1>, f32
          %11 = llvm.load %10 <alignment = 4> : !llvm.ptr<1> -> vector<1xf32>
          %12 = llvm.getelementptr %arg1[%9] : (!llvm.ptr<1>, i64) -> !llvm.ptr<1>, f32
          %13 = llvm.load %12 <alignment = 4> : !llvm.ptr<1> -> vector<1xf32>
          %14 = llvm.fmul %11, %13 : vector<1xf32>
          %15 = llvm.getelementptr %arg2[%9] : (!llvm.ptr<1>, i64) -> !llvm.ptr<1>, f32
          llvm.store %14, %15 <alignment = 4> : vector<1xf32>, !llvm.ptr<1>
          %16 = llvm.add %7, %2 : i32
          llvm.br ^bb1(%16 : i32)
        ^bb3:  // pred: ^bb1
          llvm.return
        }
      }
    }
  }
  util.func public @simple_mul(%arg0: !hal.buffer_view, %arg1: !hal.buffer_view) -> !hal.buffer_view attributes {iree.abi.stub, iree.reflection = {iree.abi.declaration = "sync func @simple_mul(%input0: tensor<4xf32>, %input1: tensor<4xf32>) -> (%output0: tensor<4xf32>)"}} {
    %c0 = arith.constant 0 : index
    %c16 = arith.constant 16 : index
    %c4 = arith.constant 4 : index
    %element_type_f32 = hal.element_type<f32> : i32
    %dense_row_major = hal.encoding_type<dense_row_major> : i32
    hal.buffer_view.assert<%arg0 : !hal.buffer_view> message("input0") shape([%c4]) type(%element_type_f32) encoding(%dense_row_major)
    %0 = stream.tensor.import on(#hal.device.affinity<@__device_0>) %arg0 : !hal.buffer_view -> tensor<4xf32> in !stream.resource<external>{%c16}
    hal.buffer_view.assert<%arg1 : !hal.buffer_view> message("input1") shape([%c4]) type(%element_type_f32) encoding(%dense_row_major)
    %1 = stream.tensor.import on(#hal.device.affinity<@__device_0>) %arg1 : !hal.buffer_view -> tensor<4xf32> in !stream.resource<external>{%c16}
    %result, %result_timepoint = stream.resource.alloca uninitialized on(#hal.device.affinity<@__device_0>) : !stream.resource<external>{%c16} => !stream.timepoint
    %2 = stream.cmd.execute on(#hal.device.affinity<@__device_0>) await(%result_timepoint) => with(%0 as %arg2: !stream.resource<external>{%c16}, %1 as %arg3: !stream.resource<external>{%c16}, %result as %arg4: !stream.resource<external>{%c16}) {
      stream.cmd.dispatch @simple_mul_dispatch_0::@cuda_nvptx_fb::@simple_mul_dispatch_0_elementwise_4_f32 {
        ro %arg2[%c0 for %c16] : !stream.resource<external>{%c16},
        ro %arg3[%c0 for %c16] : !stream.resource<external>{%c16},
        wo %arg4[%c0 for %c16] : !stream.resource<external>{%c16}
      }
    } => !stream.timepoint
    %3 = stream.timepoint.await %2 => %result : !stream.resource<external>{%c16}
    %4 = stream.tensor.export on(#hal.device.affinity<@__device_0>) %3 : tensor<4xf32> in !stream.resource<external>{%c16} -> !hal.buffer_view
    util.return %4 : !hal.buffer_view
  }
}

8. HAL: package the kernel and construct runtime execution

HAL completes the next major transition: the device implementation is serialized, and the surrounding schedule becomes explicit HAL operations.

The kernel becomes executable data

The variant’s LLVM/NVVM module is replaced by:

hal.executable.binary public @cuda_nvptx_fb attributes {
  data = dense<"0x..."> : vector<1484xi8>,
  format = "cuda-nvptx-fb",
  mime_type = "application/x-flatbuffers"
}

This 1,484-byte payload contains a CUDA executable FlatBuffer, including PTX and entry-point metadata. Decoding the artifact reveals:

.version 7.6
.target sm_60
.address_size 64

Its core instructions read %tid.x, skip threads with IDs greater than 3, load A and B, execute mul.rn.f32, and store C. The earlier loop has simplified to one bounds check because a participating thread can execute at most one iteration.

The package contains PTX rather than final GPU machine instructions. The CUDA loading path and driver handle the subsequent device executable processing.

Initialization creates reusable state

util.initializer code enumerates runtime devices, queries for CUDA and executable-format support, and creates a !hal.executable. Another initializer records and caches a command buffer through @__simple_mul_memoize_apply.

The recording uses binding slots 0, 1, and 2 and fixed workgroup counts (1, 1, 1). It also records an execution barrier and finalizes the command buffer. This barrier orders command execution; it is not a barrier among CUDA threads.

Initialization prepares executable and command objects. It does not perform the model multiplication.

Each invocation supplies resources and submits work

The public function checks the input views and underlying buffers, obtains a 16-byte output with hal.device.queue.alloca, and submits the cached command buffer through hal.device.queue.execute.indirect.

The indirect submission supplies the actual A, B, and C buffers for the recorded slots. This separates a stable command structure from per-invocation data. Therefore we can outline such dependency:

Allocate C ──signal allocation fence──► Execute cached command buffer
                                               │
                                       signal completion fence
                                               │
                                      Await completion, return C

hal.buffer_view.create wraps the result with shape [4] and type f32. Returning that view does not itself copy the result to CPU memory. This function expects compatible input buffers; any upload performed by a calling tool is outside this function’s IR.

Complete result after HAL:

module {
  util.global private @__device_0 : !hal.device
  util.initializer {
    %c18_i32 = arith.constant 18 : i32
    %false = arith.constant false
    %c0 = arith.constant 0 : index
    %c1 = arith.constant 1 : index
    %0 = util.null : !hal.device
    %device_count = hal.devices.count : index
    cf.br ^bb1(%c0, %c0, %0 : index, index, !hal.device)
  ^bb1(%1: index, %2: index, %3: !hal.device):  // 2 preds: ^bb0, ^bb4
    %4 = util.cmp.eq %3, %0 : !hal.device
    %5 = arith.cmpi slt, %1, %device_count : index
    %6 = arith.andi %4, %5 : i1
    cf.cond_br %6, ^bb2, ^bb5
  ^bb2:  // pred: ^bb1
    %device_n = hal.devices.get %1 : !hal.device
    %ok, %value = hal.device.query<%device_n : !hal.device> key("hal.device.id" :: "cuda") : i1, i1 = false
    cf.cond_br %value, ^bb3, ^bb4(%false : i1)
  ^bb3:  // pred: ^bb2
    %ok_0, %value_1 = hal.device.query<%device_n : !hal.device> key("hal.executable.format" :: "cuda-nvptx-fb") : i1, i1 = false
    cf.br ^bb4(%value_1 : i1)
  ^bb4(%7: i1):  // 2 preds: ^bb2, ^bb3
    %8 = arith.cmpi eq, %2, %c0 : index
    %9 = arith.select %7, %c1, %c0 : index
    %10 = arith.addi %2, %9 : index
    %11 = arith.andi %7, %8 : i1
    %12 = arith.select %11, %device_n, %0 : !hal.device
    %13 = arith.addi %1, %c1 : index
    cf.br ^bb1(%13, %10, %12 : index, index, !hal.device)
  ^bb5:  // pred: ^bb1
    cf.cond_br %4, ^bb6, ^bb7
  ^bb6:  // pred: ^bb5
    util.status.check_ok %c18_i32, "HAL device `__device_0` not found or unavailable: #hal.device.target<\22cuda\22, [#hal.executable.target<\22cuda\22, \22cuda-nvptx-fb\22, {iree_codegen.target_info = #iree_gpu.target<arch = \22sm_60\22, features = \22+ptx76\22, wgp = <compute = fp64|fp32|fp16|int64|int32|int16|int8, storage = b64|b32|b16|b8, subgroup = shuffle|arithmetic, subgroup_size_choices = [32], max_workgroup_sizes = [1024, 1024, 1024], max_thread_count_per_workgroup = 1024, max_workgroup_memory_bytes = 49152, max_workgroup_counts = [2147483647, 65535, 65535]>>}>]>"
    cf.br ^bb7
  ^bb7:  // 2 preds: ^bb5, ^bb6
    util.global.store %3, @__device_0 : !hal.device
    util.return
  }
  util.global private @__device_0_query_0_hal_executable_format_cuda_nvptx_fb : i1
  util.initializer {
    %__device_0 = util.global.load @__device_0 : !hal.device
    %ok, %value = hal.device.query<%__device_0 : !hal.device> key("hal.executable.format" :: "cuda-nvptx-fb") : i1, i1 = false
    util.global.store %value, @__device_0_query_0_hal_executable_format_cuda_nvptx_fb : i1
    util.return
  }
  util.global private @__device_0_executable_0_simple_mul_dispatch_0 : !hal.executable
  util.initializer {
    %c-1_i64 = arith.constant -1 : i64
    %c-1 = arith.constant -1 : index
    %c0 = arith.constant 0 : index
    %c14_i32 = arith.constant 14 : i32
    %0 = util.null : !hal.executable
    %__device_0_query_0_hal_executable_format_cuda_nvptx_fb = util.global.load @__device_0_query_0_hal_executable_format_cuda_nvptx_fb : i1
    %__device_0 = util.global.load @__device_0 : !hal.device
    %1 = arith.select %__device_0_query_0_hal_executable_format_cuda_nvptx_fb, %c0, %c-1 : index
    %2 = arith.cmpi eq, %1, %c0 : index
    cf.cond_br %2, ^bb1, ^bb2
  ^bb1:  // pred: ^bb0
    %executable = hal.executable.create device(%__device_0 : !hal.device) affinity(%c-1_i64) target(@simple_mul_dispatch_0::@cuda_nvptx_fb) : !hal.executable
    cf.br ^bb3(%executable : !hal.executable)
  ^bb2:  // pred: ^bb0
    util.status.check_ok %c14_i32, "HAL device `__device_0` does not support any variant of executable `simple_mul_dispatch_0`; available formats: [cuda-nvptx-fb]"
    cf.br ^bb3(%0 : !hal.executable)
  ^bb3(%3: !hal.executable):  // 2 preds: ^bb1, ^bb2
    util.global.store %3, @__device_0_executable_0_simple_mul_dispatch_0 : !hal.executable
    util.return
  }
  hal.executable private @simple_mul_dispatch_0 {
    hal.executable.binary public @cuda_nvptx_fb attributes {data = dense<"..."> : vector<1484xi8>, format = "cuda-nvptx-fb", mime_type = "application/x-flatbuffers"}
  }
  util.func private @__simple_mul_memoize_apply() -> !hal.command_buffer attributes {inlining_policy = #util.inline.never} {
    %c-1_i64 = arith.constant -1 : i64
    %c3 = arith.constant 3 : index
    %c0 = arith.constant 0 : index
    %c1 = arith.constant 1 : index
    %c2 = arith.constant 2 : index
    %c16 = arith.constant 16 : index
    %__device_0 = util.global.load immutable @__device_0 : !hal.device
    %__device_0_executable_0_simple_mul_dispatch_0 = util.global.load immutable @__device_0_executable_0_simple_mul_dispatch_0 : !hal.executable
    %cmd = hal.command_buffer.create device(%__device_0 : !hal.device) mode("None") categories("Transfer|Dispatch") affinity(%c-1_i64) bindings(%c3) : !hal.command_buffer
    %function_id = hal.executable.lookup.function target(%__device_0_executable_0_simple_mul_dispatch_0 : !hal.executable) function(@simple_mul_dispatch_0::@cuda_nvptx_fb::@simple_mul_dispatch_0_elementwise_4_f32) : i64
    hal.command_buffer.dispatch<%cmd : !hal.command_buffer> target(%__device_0_executable_0_simple_mul_dispatch_0 : !hal.executable)[%function_id] workgroups([%c1, %c1, %c1]) bindings([
      (%c0 : index)[%c0, %c16], 
      (%c1 : index)[%c0, %c16], 
      (%c2 : index)[%c0, %c16]
    ]) flags("None")
    hal.command_buffer.execution_barrier<%cmd : !hal.command_buffer> source("Dispatch|Transfer|CommandRetire") target("CommandIssue|Dispatch|Transfer") flags("None")
    hal.command_buffer.finalize<%cmd : !hal.command_buffer>
    util.return %cmd : !hal.command_buffer
  }
  util.global private @__simple_mul_memoize_result_0_device_0 : !hal.command_buffer
  util.initializer {
    %0 = util.call @__simple_mul_memoize_apply() : () -> !hal.command_buffer
    util.global.store %0, @__simple_mul_memoize_result_0_device_0 : !hal.command_buffer
    util.return
  }
  util.func public @simple_mul(%arg0: !hal.buffer_view, %arg1: !hal.buffer_view) -> !hal.buffer_view attributes {iree.abi.stub, iree.reflection = {iree.abi.declaration = "sync func @simple_mul(%input0: tensor<4xf32>, %input1: tensor<4xf32>) -> (%output0: tensor<4xf32>)"}} {
    %c-1_i32 = arith.constant -1 : i32
    %c0_i64 = arith.constant 0 : i64
    %0 = util.null : !hal.fence
    %c-1_i64 = arith.constant -1 : i64
    %c0 = arith.constant 0 : index
    %c16 = arith.constant 16 : index
    %c4 = arith.constant 4 : index
    %memory_type = hal.memory_type<"DeviceVisible|DeviceLocal"> : i32
    %buffer_usage = hal.buffer_usage<"TransferSource|TransferTarget|Transfer|DispatchStorageRead|DispatchStorageWrite|DispatchStorage"> : i32
    %__device_0 = util.global.load immutable @__device_0 : !hal.device
    %__simple_mul_memoize_result_0_device_0 = util.global.load immutable @__simple_mul_memoize_result_0_device_0 : !hal.command_buffer
    %element_type_f32 = hal.element_type<f32> : i32
    %dense_row_major = hal.encoding_type<dense_row_major> : i32
    hal.buffer_view.assert<%arg0 : !hal.buffer_view> message("input0") shape([%c4]) type(%element_type_f32) encoding(%dense_row_major)
    %buffer = hal.buffer_view.buffer<%arg0 : !hal.buffer_view> : !hal.buffer
    %allocator = hal.device.allocator<%__device_0 : !hal.device> : !hal.allocator
    hal.buffer.assert<%buffer : !hal.buffer> message("tensor") allocator(%allocator : !hal.allocator) minimum_length(%c16) type(DeviceVisible) usage("TransferSource|TransferTarget|Transfer|DispatchStorageRead|DispatchStorageWrite|DispatchStorage")
    hal.buffer_view.assert<%arg1 : !hal.buffer_view> message("input1") shape([%c4]) type(%element_type_f32) encoding(%dense_row_major)
    %buffer_0 = hal.buffer_view.buffer<%arg1 : !hal.buffer_view> : !hal.buffer
    hal.buffer.assert<%buffer_0 : !hal.buffer> message("tensor") allocator(%allocator : !hal.allocator) minimum_length(%c16) type(DeviceVisible) usage("TransferSource|TransferTarget|Transfer|DispatchStorageRead|DispatchStorageWrite|DispatchStorage")
    %fence = hal.fence.create device(%__device_0 : !hal.device) flags("None") : !hal.fence
    %transient_buffer = hal.device.queue.alloca<%__device_0 : !hal.device> affinity(%c-1_i64) wait(%0) signal(%fence) pool(%c0_i64) type(%memory_type) usage(%buffer_usage) flags("None") : !hal.buffer{%c16}
    %fence_1 = hal.fence.create device(%__device_0 : !hal.device) flags("None") : !hal.fence
    hal.device.queue.execute.indirect<%__device_0 : !hal.device> affinity(%c-1_i64) wait(%fence) signal(%fence_1) commands(%__simple_mul_memoize_result_0_device_0) bindings([
      (%buffer : !hal.buffer)[%c0, %c16], 
      (%buffer_0 : !hal.buffer)[%c0, %c16], 
      (%transient_buffer : !hal.buffer)[%c0, %c16]
    ]) flags("None")
    %status = hal.fence.await until([%fence_1]) timeout_millis(%c-1_i32) flags("None") : i32
    util.status.check_ok %status, "failed to wait on timepoint"
    %view = hal.buffer_view.create buffer(%transient_buffer : !hal.buffer)[%c0, %c16] shape([%c4]) type(%element_type_f32) encoding(%dense_row_major) : !hal.buffer_view
    util.return %view : !hal.buffer_view
  }
}

9. VM: encode orchestration as portable control code

Finally, the VM pipeline creates a vm.module with functions, imports, exports, globals, and read-only data.

HAL objects become VM references:

!hal.buffer_view → !vm.ref<!hal.buffer_view>
!hal.device      → !vm.ref<!hal.device>

A HAL operation becomes a call to a runtime import:

vm.import private @hal.buffer_view.buffer(
    %view: !vm.ref<!hal.buffer_view>) -> !vm.ref<!hal.buffer>

%buffer = vm.call @hal.buffer_view.buffer(%view)
    : (!vm.ref<!hal.buffer_view>) -> !vm.ref<!hal.buffer>

This is the runtime boundary: the VM executes control instructions and invokes the HAL module, whose implementation manages the device resources. Import minimum_version attributes constrain the HAL module interface version, not the LLVM revision. Strings and executable bytes become vm.rodata. The CUDA payload remains 1,484 bytes. vm.const.ref.rodata supplies references to it when initialization calls @hal.executable.create.

Waiting can suspend VM execution

The fence wait becomes:

vm.call.variadic.yieldable @hal.fence.await(...)
    -> ^resume(i32)

If the import completes immediately, control proceeds to the successor with its result. If it yields, execution resumes there after waiting. The successor checks the returned status, creates the output view on success, or executes vm.fail after cleanup on failure.

The synchronous public ABI requires completion before returning a result. Internally, the VM can support suspension and resumption; the HAL runtime may alternatively block a native thread in its synchronous mode. A synchronous application contract does not require every layer to implement waiting identically.

module attributes {vm.toplevel} {
  vm.module public @module {
    vm.global.ref private mutable @__device_0 : !vm.ref<!hal.device>
    vm.global.ref private mutable @__device_0_executable_0_simple_mul_dispatch_0 : !vm.ref<!hal.executable>
    vm.global.ref private mutable @__simple_mul_memoize_result_0_device_0 : !vm.ref<!hal.command_buffer>
    vm.rodata private @_utf8_hal_device_id_C6650FF277232B5A {alignment = 1 : i64} "hal.device.id"
    vm.rodata private @_utf8_cuda_342A062BE26F4639 {alignment = 1 : i64} "cuda"
    vm.rodata private @_utf8_hal_executable_format_E03EECB63A2AAF52 {alignment = 1 : i64} "hal.executable.format"
    vm.rodata private @_utf8_cuda_nvptx_fb_91E5BBD19AE47511 {alignment = 1 : i64} "cuda-nvptx-fb"
    vm.rodata private @simple_mul_dispatch_0_cuda_nvptx_fb {alignment = 16 : i64, mime_type = "application/x-flatbuffers"} dense<"0x43444131000000008C0500000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000008000000434441319AFAFFFF08000000100100000100000004000000A6FAFFFF640000002000000001000000010000000C0000007C00000020000000030000000300000000000000030000000000000002000000000000002700000073696D706C655F6D756C5F64697370617463685F305F656C656D656E74776973655F345F663332002700000073696D706C655F6D756C5F64697370617463685F305F656C656D656E74776973655F345F6633320042FBFFFF40000000040000004EFBFFFF08000000010000002600000073616D706C65732F7374617469635F6C6962726172792F73696D706C655F6D756C2E6D6C697200002700000073696D706C655F6D756C5F64697370617463685F305F656C656D656E74776973655F345F663332000100000004000000C0FBFFFF04000000330400002F2F0A2F2F2047656E657261746564206279204C4C564D204E56505458204261636B2D456E640A2F2F0A0A2E76657273696F6E20372E360A2E74617267657420736D5F36300A2E616464726573735F73697A652036340A0A092F2F202E676C6F626C0973696D706C655F6D756C5F64697370617463685F305F656C656D656E74776973655F345F6633320A0A2E76697369626C65202E656E7472792073696D706C655F6D756C5F64697370617463685F305F656C656D656E74776973655F345F663332280A092E706172616D202E753634202E707472202E676C6F62616C202E616C69676E2031362073696D706C655F6D756C5F64697370617463685F305F656C656D656E74776973655F345F6633325F706172616D5F302C0A092E706172616D202E753634202E707472202E676C6F62616C202E616C69676E2031362073696D706C655F6D756C5F64697370617463685F305F656C656D656E74776973655F345F6633325F706172616D5F312C0A092E706172616D202E753634202E707472202E676C6F62616C202E616C69676E2031362073696D706C655F6D756C5F64697370617463685F305F656C656D656E74776973655F345F6633325F706172616D5F320A290A2E6D61786E7469642033322C20312C20310A7B0A092E726567202E70726564200925703C323E3B0A092E726567202E623332200925723C353E3B0A092E726567202E62363420092572643C383E3B0A0A096D6F762E75333220092572312C20257469642E783B0A09736574702E67742E75333220092570312C202572312C20333B0A0940257031206272612009244C5F5F4242305F323B0A096C642E706172616D2E6236342009257264342C205B73696D706C655F6D756C5F64697370617463685F305F656C656D656E74776973655F345F6633325F706172616D5F305D3B0A096C642E706172616D2E6236342009257264352C205B73696D706C655F6D756C5F64697370617463685F305F656C656D656E74776973655F345F6633325F706172616D5F315D3B0A096C642E706172616D2E6236342009257264362C205B73696D706C655F6D756C5F64697370617463685F305F656C656D656E74776973655F345F6633325F706172616D5F325D3B0A096D756C2E776964652E7533322009257264372C202572312C20343B0A096164642E7336342009257264312C20257264342C20257264373B0A096164642E7336342009257264322C20257264352C20257264373B0A096164642E7336342009257264332C20257264362C20257264373B0A096C642E676C6F62616C2E6E632E62333220092572322C205B257264315D3B0A096C642E676C6F62616C2E6E632E62333220092572332C205B257264325D3B0A096D756C2E726E2E66333220092572342C202572322C202572333B0A0973742E676C6F62616C2E62333220095B257264335D2C202572343B0A244C5F5F4242305F323A0A097265743B0A0A7D0A0006000800040008000C00040008001400200000000400080000000000140018001C000000"> : vector<1484xi8>
    vm.rodata private @_utf8_simple_mul_dispatch_0_elem_BAECB3686C800B2E {alignment = 1 : i64} "simple_mul_dispatch_0_elementwise_4_f32"
    vm.func private @__simple_mul_memoize_apply() -> !vm.ref<!hal.command_buffer> attributes {vm.unwind, inlining_policy = #util.inline.never} {
      %c13 = vm.const.i32 13
      %c28 = vm.const.i32 28
      %c2 = vm.const.i32 2
      %null = vm.const.ref.zero : !vm.ref<!hal.buffer>
      %c1 = vm.const.i32 1
      %c3 = vm.const.i32 3
      %zero = vm.const.i32.zero
      %c16 = vm.const.i64 16
      %zero_0 = vm.const.i64.zero
      %c-1 = vm.const.i64 -1
      %__device_0 = vm.global.load.ref @__device_0 : !vm.ref<!hal.device>
      %__device_0_executable_0_simple_mul_dispatch_0 = vm.global.load.ref @__device_0_executable_0_simple_mul_dispatch_0 : !vm.ref<!hal.executable>
      %ref = vm.call @hal.command_buffer.create(%__device_0, %zero, %c3, %c-1, %c3) : (!vm.ref<!hal.device>, i32, i32, i64, i32) -> !vm.ref<!hal.command_buffer>
      %_utf8_simple_mul_dispatch_0_elem_BAECB3686C800B2E = vm.const.ref.rodata @_utf8_simple_mul_dispatch_0_elem_BAECB3686C800B2E : !vm.buffer
      %0 = vm.call @hal.executable.lookup.function(%__device_0_executable_0_simple_mul_dispatch_0, %_utf8_simple_mul_dispatch_0_elem_BAECB3686C800B2E) {nosideeffects} : (!vm.ref<!hal.executable>, !vm.buffer) -> i64
      vm.call.variadic @hal.command_buffer.dispatch(%ref, %__device_0_executable_0_simple_mul_dispatch_0, %0, %c1, %c1, %c1, %zero_0, [], [(%zero, %zero, %null, %zero_0, %c16), (%zero, %c1, %null, %zero_0, %c16), (%zero, %c2, %null, %zero_0, %c16)]) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable>, i64, i32, i32, i32, i64, i32 ..., tuple<i32, i32, !vm.ref<!hal.buffer>, i64, i64> ...)
      vm.call @hal.command_buffer.execution_barrier(%ref, %c28, %c13, %zero_0) : (!vm.ref<!hal.command_buffer>, i32, i32, i64) -> ()
      vm.call @hal.command_buffer.finalize(%ref) : (!vm.ref<!hal.command_buffer>) -> ()
      vm.return %ref : !vm.ref<!hal.command_buffer>
    }
    vm.import private @hal.buffer.assert(%buffer : !vm.ref<!hal.buffer>, %message : !vm.buffer, %allocator : !vm.ref<!hal.allocator>, %minimum_length : i64, %memory_types : i32, %buffer_usage : i32)
    vm.import private @hal.buffer_view.create(%buffer : !vm.ref<!hal.buffer>, %source_offset : i64, %source_length : i64, %element_type : i32, %encoding_type : i32, %shape : i64 ...) -> !vm.ref<!hal.buffer_view> attributes {nosideeffects}
    vm.import private @hal.buffer_view.assert(%buffer_view : !vm.ref<!hal.buffer_view>, %message : !vm.buffer, %element_type : i32, %encoding_type : i32, %shape : i64 ...)
    vm.import private @hal.buffer_view.buffer(%buffer_view : !vm.ref<!hal.buffer_view>) -> !vm.ref<!hal.buffer> attributes {nosideeffects}
    vm.import private @hal.command_buffer.create(%device : !vm.ref<!hal.device>, %modes : i32, %command_categories : i32, %queue_affinity : i64, %binding_capacity : i32) -> !vm.ref<!hal.command_buffer> attributes {minimum_version = 6 : i32}
    vm.import private @hal.command_buffer.finalize(%command_buffer : !vm.ref<!hal.command_buffer>)
    vm.import private @hal.command_buffer.execution_barrier(%command_buffer : !vm.ref<!hal.command_buffer>, %source_stage_mask : i32, %target_stage_mask : i32, %flags : i64)
    vm.import private @hal.command_buffer.dispatch(%command_buffer : !vm.ref<!hal.command_buffer>, %executable : !vm.ref<!hal.executable>, %function_id : i64, %workgroup_x : i32, %workgroup_y : i32, %workgroup_z : i32, %flags : i64, %constants : i32 ..., %bindings : tuple<i32, i32, !vm.ref<!hal.buffer>, i64, i64> ...) attributes {minimum_version = 7 : i32}
    vm.import private @hal.device.allocator(%device : !vm.ref<!hal.device>) -> !vm.ref<!hal.allocator> attributes {nosideeffects}
    vm.import private @hal.device.query.i64(%device : !vm.ref<!hal.device>, %category : !vm.buffer, %key : !vm.buffer) -> (i32, i64) attributes {nosideeffects}
    vm.import private @hal.device.queue.alloca(%device : !vm.ref<!hal.device>, %queue_affinity : i64, %wait_fence : !vm.ref<!hal.fence>, %signal_fence : !vm.ref<!hal.fence>, %pool : i64, %memory_types : i32, %buffer_usage : i32, %allocation_size : i64, %flags : i64) -> !vm.ref<!hal.buffer>
    vm.import private @hal.device.queue.execute.indirect(%device : !vm.ref<!hal.device>, %queue_affinity : i64, %wait_fence : !vm.ref<!hal.fence>, %signal_fence : !vm.ref<!hal.fence>, %command_buffer : !vm.ref<!hal.command_buffer>, %flags : i64, %binding_table : tuple<!vm.ref<!hal.buffer>, i64, i64> ...)
    vm.import private @hal.devices.count() -> i32 attributes {nosideeffects}
    vm.import private @hal.devices.get(%index : i32) -> !vm.ref<!hal.device> attributes {nosideeffects}
    vm.import private @hal.executable.create(%device : !vm.ref<!hal.device>, %queue_affinity : i64, %executable_format : !vm.buffer, %executable_data : !vm.buffer, %constants : !vm.buffer) -> !vm.ref<!hal.executable> attributes {nosideeffects}
    vm.import private @hal.executable.lookup.function(%executable : !vm.ref<!hal.executable>, %function_name : !vm.buffer) -> i64 attributes {nosideeffects, minimum_version = 7 : i32}
    vm.import private @hal.fence.create(%device : !vm.ref<!hal.device>, %flags : i64) -> !vm.ref<!hal.fence>
    vm.import private @hal.fence.await(%timeout_millis : i32, %flags : i64, %fences : !vm.ref<!hal.fence> ...) -> i32 attributes {vm.yield}
    vm.rodata private @_utf8_input0_DCE99660CEB3F6B {alignment = 1 : i64} "input0"
    vm.rodata private @_utf8_tensor_FC1814BC4A58F22A {alignment = 1 : i64} "tensor"
    vm.rodata private @_utf8_input1_B898B726583C85DA {alignment = 1 : i64} "input1"
    vm.func private @simple_mul(%arg0: !vm.ref<!hal.buffer_view>, %arg1: !vm.ref<!hal.buffer_view>) -> !vm.ref<!hal.buffer_view> attributes {iree.reflection = {iree.abi.declaration = "sync func @simple_mul(%input0: tensor<4xf32>, %input1: tensor<4xf32>) -> (%output0: tensor<4xf32>)"}, vm.unwind, vm.yield} {
      %c16 = vm.const.i32 16
      %c1 = vm.const.i32 1
      %c553648160 = vm.const.i32 553648160
      %c3075 = vm.const.i32 3075
      %c48 = vm.const.i32 48
      %c4 = vm.const.i64 4
      %c16_0 = vm.const.i64 16
      %zero = vm.const.i64.zero
      %c-1 = vm.const.i64 -1
      %null = vm.const.ref.zero : !vm.ref<!hal.fence>
      %c-1_1 = vm.const.i32 -1
      %__device_0 = vm.global.load.ref @__device_0 : !vm.ref<!hal.device>
      %__simple_mul_memoize_result_0_device_0 = vm.global.load.ref @__simple_mul_memoize_result_0_device_0 : !vm.ref<!hal.command_buffer>
      %_utf8_input0_DCE99660CEB3F6B = vm.const.ref.rodata @_utf8_input0_DCE99660CEB3F6B : !vm.buffer
      vm.call.variadic @hal.buffer_view.assert(%arg0, %_utf8_input0_DCE99660CEB3F6B, %c553648160, %c1, [%c4]) : (!vm.ref<!hal.buffer_view>, !vm.buffer, i32, i32, i64 ...)
      %ref = vm.call @hal.buffer_view.buffer(%arg0) {nosideeffects} : (!vm.ref<!hal.buffer_view>) -> !vm.ref<!hal.buffer>
      %ref_2 = vm.call @hal.device.allocator(%__device_0) {nosideeffects} : (!vm.ref<!hal.device>) -> !vm.ref<!hal.allocator>
      %_utf8_tensor_FC1814BC4A58F22A = vm.const.ref.rodata @_utf8_tensor_FC1814BC4A58F22A : !vm.buffer
      vm.call @hal.buffer.assert(%ref, %_utf8_tensor_FC1814BC4A58F22A, %ref_2, %c16_0, %c16, %c3075) : (!vm.ref<!hal.buffer>, !vm.buffer, !vm.ref<!hal.allocator>, i64, i32, i32) -> ()
      %_utf8_input1_B898B726583C85DA = vm.const.ref.rodata @_utf8_input1_B898B726583C85DA : !vm.buffer
      vm.call.variadic @hal.buffer_view.assert(%arg1, %_utf8_input1_B898B726583C85DA, %c553648160, %c1, [%c4]) : (!vm.ref<!hal.buffer_view>, !vm.buffer, i32, i32, i64 ...)
      %ref_3 = vm.call @hal.buffer_view.buffer(%arg1) {nosideeffects} : (!vm.ref<!hal.buffer_view>) -> !vm.ref<!hal.buffer>
      vm.call @hal.buffer.assert(%ref_3, %_utf8_tensor_FC1814BC4A58F22A, %ref_2, %c16_0, %c16, %c3075) : (!vm.ref<!hal.buffer>, !vm.buffer, !vm.ref<!hal.allocator>, i64, i32, i32) -> ()
      %ref_4 = vm.call @hal.fence.create(%__device_0, %zero) : (!vm.ref<!hal.device>, i64) -> !vm.ref<!hal.fence>
      %ref_5 = vm.call @hal.device.queue.alloca(%__device_0, %c-1, %null, %ref_4, %zero, %c48, %c3075, %c16_0, %zero) : (!vm.ref<!hal.device>, i64, !vm.ref<!hal.fence>, !vm.ref<!hal.fence>, i64, i32, i32, i64, i64) -> !vm.ref<!hal.buffer>
      %ref_6 = vm.call @hal.fence.create(%__device_0, %zero) : (!vm.ref<!hal.device>, i64) -> !vm.ref<!hal.fence>
      vm.call.variadic @hal.device.queue.execute.indirect(%__device_0, %c-1, %ref_4, %ref_6, %__simple_mul_memoize_result_0_device_0, %zero, [(%ref, %zero, %c16_0), (%ref_3, %zero, %c16_0), (%ref_5, %zero, %c16_0)]) : (!vm.ref<!hal.device>, i64, !vm.ref<!hal.fence>, !vm.ref<!hal.fence>, !vm.ref<!hal.command_buffer>, i64, tuple<!vm.ref<!hal.buffer>, i64, i64> ...)
      vm.call.variadic.yieldable @hal.fence.await(%c-1_1, %zero, %ref_6) {segment_sizes = dense<[-1, -1, 1]> : vector<3xi16>, segment_types = [i32, i64, !vm.ref<!hal.fence>]} : (i32, i64, !vm.ref<!hal.fence>) -> ^bb1 (i32)
    ^bb1(%0: i32):  // pred: ^bb0
      vm.cond_br %0, ^bb3, ^bb2
    ^bb2:  // pred: ^bb1
      %ref_7 = vm.call.variadic @hal.buffer_view.create(%ref_5, %zero, %c16_0, %c553648160, %c1, [%c4]) {nosideeffects} : (!vm.ref<!hal.buffer>, i64, i64, i32, i32, i64 ...) -> !vm.ref<!hal.buffer_view>
      vm.return %ref_7 : !vm.ref<!hal.buffer_view>
    ^bb3:  // pred: ^bb1
      vm.discard.refs %ref_5 : !vm.ref<!hal.buffer>
      vm.fail %0, "failed to wait on timepoint"
    }
    vm.export @simple_mul attributes {iree.abi.stub, iree.reflection = {iree.abi.declaration = "sync func @simple_mul(%input0: tensor<4xf32>, %input1: tensor<4xf32>) -> (%output0: tensor<4xf32>)"}}
    vm.export @__init
    vm.func private @__init() attributes {vm.unwind} {
      %c1 = vm.const.i32 1
      %null = vm.const.ref.zero : !vm.buffer
      %c14 = vm.const.i32 14
      %c-1 = vm.const.i64 -1
      %c18 = vm.const.i32 18
      %zero = vm.const.i32.zero
      %zero_0 = vm.const.i64.zero
      %c1_1 = vm.const.i64 1
      %null_2 = vm.const.ref.zero : !vm.ref<!hal.device>
      %0 = vm.call @hal.devices.count() {nosideeffects} : () -> i32
      %1 = vm.ext.i32.i64.s %0 : i32 -> i64
      vm.br ^bb1(%zero_0, %zero_0, %null_2 : i64, i64, !vm.ref<!hal.device>)
    ^bb1(%2: i64, %3: i64, %4: !vm.ref<!hal.device>):  // 2 preds: ^bb0, ^bb4
      %rnz = vm.cmp.nz.ref %4 : !vm.ref<!hal.device>
      %5 = vm.xor.i32 %rnz, %c1 : i32
      %slt = vm.cmp.lt.i64.s %2, %1 : i64
      %6 = vm.and.i32 %5, %slt : i32
      vm.cond_br %6, ^bb2, ^bb5
    ^bb2:  // pred: ^bb1
      vm.discard.refs %4 : !vm.ref<!hal.device>
      %7 = vm.trunc.i64.i32 %2 : i64 -> i32
      %ref = vm.call @hal.devices.get(%7) {nosideeffects} : (i32) -> !vm.ref<!hal.device>
      %_utf8_hal_device_id_C6650FF277232B5A = vm.const.ref.rodata @_utf8_hal_device_id_C6650FF277232B5A : !vm.buffer
      %_utf8_cuda_342A062BE26F4639 = vm.const.ref.rodata @_utf8_cuda_342A062BE26F4639 : !vm.buffer
      %8:2 = vm.call @hal.device.query.i64(%ref, %_utf8_hal_device_id_C6650FF277232B5A, %_utf8_cuda_342A062BE26F4639) {nosideeffects} : (!vm.ref<!hal.device>, !vm.buffer, !vm.buffer) -> (i32, i64)
      %nz = vm.cmp.nz.i64 %8#1 : i64
      %9 = vm.select.i32 %8#0, %nz, %zero : i32
      vm.cond_br %9, ^bb3, ^bb4(%zero : i32)
    ^bb3:  // pred: ^bb2
      %_utf8_hal_executable_format_E03EECB63A2AAF52 = vm.const.ref.rodata @_utf8_hal_executable_format_E03EECB63A2AAF52 : !vm.buffer
      %_utf8_cuda_nvptx_fb_91E5BBD19AE47511 = vm.const.ref.rodata @_utf8_cuda_nvptx_fb_91E5BBD19AE47511 : !vm.buffer
      %10:2 = vm.call @hal.device.query.i64(%ref, %_utf8_hal_executable_format_E03EECB63A2AAF52, %_utf8_cuda_nvptx_fb_91E5BBD19AE47511) {nosideeffects} : (!vm.ref<!hal.device>, !vm.buffer, !vm.buffer) -> (i32, i64)
      %nz_3 = vm.cmp.nz.i64 %10#1 : i64
      %11 = vm.select.i32 %10#0, %nz_3, %zero : i32
      vm.br ^bb4(%11 : i32)
    ^bb4(%12: i32):  // 2 preds: ^bb2, ^bb3
      %eq = vm.cmp.eq.i64 %3, %zero_0 : i64
      %13 = vm.select.i64 %12, %c1_1, %zero_0 : i64
      %14 = vm.add.i64 %3, %13 : i64
      %15 = vm.and.i32 %12, %eq : i32
      %ref_4 = vm.select.ref %15, %ref, %null_2 : !vm.ref<!hal.device>
      %16 = vm.add.i64 %2, %c1_1 : i64
      vm.br ^bb1(%16, %14, %ref_4 : i64, i64, !vm.ref<!hal.device>)
    ^bb5:  // pred: ^bb1
      vm.discard.refs %null_2 : !vm.ref<!hal.device>
      vm.cond_br %5, ^bb6, ^bb7
    ^bb6:  // pred: ^bb5
      vm.discard.refs %null, %4 : !vm.buffer, !vm.ref<!hal.device>
      vm.fail %c18, "HAL device `__device_0` not found or unavailable: #hal.device.target<\22cuda\22, [#hal.executable.target<\22cuda\22, \22cuda-nvptx-fb\22, {iree_codegen.target_info = #iree_gpu.target<arch = \22sm_60\22, features = \22+ptx76\22, wgp = <compute = fp64|fp32|fp16|int64|int32|int16|int8, storage = b64|b32|b16|b8, subgroup = shuffle|arithmetic, subgroup_size_choices = [32], max_workgroup_sizes = [1024, 1024, 1024], max_thread_count_per_workgroup = 1024, max_workgroup_memory_bytes = 49152, max_workgroup_counts = [2147483647, 65535, 65535]>>}>]>"
    ^bb7:  // pred: ^bb5
      %_utf8_hal_executable_format_E03EECB63A2AAF52_5 = vm.const.ref.rodata @_utf8_hal_executable_format_E03EECB63A2AAF52 : !vm.buffer
      %_utf8_cuda_nvptx_fb_91E5BBD19AE47511_6 = vm.const.ref.rodata @_utf8_cuda_nvptx_fb_91E5BBD19AE47511 : !vm.buffer
      %17:2 = vm.call @hal.device.query.i64(%4, %_utf8_hal_executable_format_E03EECB63A2AAF52_5, %_utf8_cuda_nvptx_fb_91E5BBD19AE47511_6) {nosideeffects} : (!vm.ref<!hal.device>, !vm.buffer, !vm.buffer) -> (i32, i64)
      %nz_7 = vm.cmp.nz.i64 %17#1 : i64
      %18 = vm.select.i32 %17#0, %nz_7, %zero : i32
      %19 = vm.select.i64 %18, %zero_0, %c-1 : i64
      %eq_8 = vm.cmp.eq.i64 %19, %zero_0 : i64
      vm.global.store.ref %4, @__device_0 : !vm.ref<!hal.device>
      vm.cond_br %eq_8, ^bb8, ^bb9
    ^bb8:  // pred: ^bb7
      %simple_mul_dispatch_0_cuda_nvptx_fb = vm.const.ref.rodata @simple_mul_dispatch_0_cuda_nvptx_fb : !vm.buffer
      %ref_9 = vm.call @hal.executable.create(%4, %c-1, %_utf8_cuda_nvptx_fb_91E5BBD19AE47511_6, %simple_mul_dispatch_0_cuda_nvptx_fb, %null) {nosideeffects} : (!vm.ref<!hal.device>, i64, !vm.buffer, !vm.buffer, !vm.buffer) -> !vm.ref<!hal.executable>
      vm.global.store.ref %ref_9, @__device_0_executable_0_simple_mul_dispatch_0 : !vm.ref<!hal.executable>
      %ref_10 = vm.call @__simple_mul_memoize_apply() : () -> !vm.ref<!hal.command_buffer>
      vm.global.store.ref %ref_10, @__simple_mul_memoize_result_0_device_0 : !vm.ref<!hal.command_buffer>
      vm.return
    ^bb9:  // pred: ^bb7
      vm.discard.refs %null, %4, %_utf8_cuda_nvptx_fb_91E5BBD19AE47511_6 : !vm.buffer, !vm.ref<!hal.device>, !vm.buffer
      vm.fail %c14, "HAL device `__device_0` does not support any variant of executable `simple_mul_dispatch_0`; available formats: [cuda-nvptx-fb]"
    }
  }
}