// Copyright 2026 The IREE Authors
//
// Licensed under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception

#include "iree/hal/replay/execute.h"

#include <errno.h>
#include <inttypes.h>
#include <stddef.h>
#include <string.h>

#if IREE_FILE_IO_ENABLE && \
    (defined(IREE_PLATFORM_ANDROID) || defined(IREE_PLATFORM_LINUX))
#include <sys/stat.h>
#include <unistd.h>
#endif  // IREE_FILE_IO_ENABLE && (IREE_PLATFORM_ANDROID ||
        // IREE_PLATFORM_LINUX)

#include "iree/hal/replay/digest.h"
#include "iree/hal/replay/file_reader.h"
#include "iree/hal/replay/format.h"
#include "iree/io/file_handle.h"

#define IREE_HAL_REPLAY_EXECUTABLE_FORMAT_CAPACITY 128
#define IREE_HAL_REPLAY_INLINE_SEMAPHORE_LIST_CAPACITY 8
#define IREE_HAL_REPLAY_INLINE_BUFFER_REF_LIST_CAPACITY 8
#define IREE_HAL_REPLAY_INLINE_BUFFER_BINDING_TABLE_CAPACITY 8
#define IREE_HAL_REPLAY_INLINE_MEMORY_BARRIER_LIST_CAPACITY 8
#define IREE_HAL_REPLAY_INLINE_BUFFER_BARRIER_LIST_CAPACITY 8

typedef struct iree_hal_replay_executable_entry_t {
  // Retained HAL executable handle used for replay dispatches.
  iree_hal_executable_t* handle;
  // Number of captured function ordinals in |function_map|.
  iree_host_size_t function_map_count;
  // Captured function ordinal to live function token mapping.
  iree_hal_executable_function_t* function_map;
} iree_hal_replay_executable_entry_t;

typedef struct iree_hal_replay_object_entry_t {
  // Captured object type stored in this entry.
  iree_hal_replay_object_type_t type;
  // Retained HAL resource pointer for the captured object.
  union {
    iree_hal_device_t* device;
    iree_hal_allocator_t* allocator;
    iree_hal_buffer_t* buffer;
    iree_hal_command_buffer_t* command_buffer;
    iree_hal_executable_cache_t* executable_cache;
    iree_hal_replay_executable_entry_t executable;
    iree_hal_semaphore_t* semaphore;
    iree_hal_event_t* event;
    iree_hal_file_t* file;
  } value;
} iree_hal_replay_object_entry_t;

typedef enum iree_hal_replay_plan_record_kind_e {
  IREE_HAL_REPLAY_PLAN_RECORD_KIND_GENERIC = 0u,
  IREE_HAL_REPLAY_PLAN_RECORD_KIND_SKIP = 1u,
  IREE_HAL_REPLAY_PLAN_RECORD_KIND_SCOPE = 2u,
  IREE_HAL_REPLAY_PLAN_RECORD_KIND_COMMAND_BUFFER_EXECUTION_BARRIER = 3u,
  IREE_HAL_REPLAY_PLAN_RECORD_KIND_COMMAND_BUFFER_DISPATCH = 4u,
  IREE_HAL_REPLAY_PLAN_RECORD_KIND_QUEUE_EXECUTE = 5u,
} iree_hal_replay_plan_record_kind_t;

typedef struct iree_hal_replay_plan_scope_t {
  // Type of scope event to publish during execution.
  iree_hal_replay_scope_event_type_t event_type;
  // Scope name borrowed from the replay file.
  iree_string_view_t name;
} iree_hal_replay_plan_scope_t;

typedef struct iree_hal_replay_plan_command_buffer_execution_barrier_t {
  // Captured command buffer object id.
  iree_hal_replay_object_id_t command_buffer_id;
  // Source execution stages for the barrier.
  iree_hal_execution_stage_t source_stage_mask;
  // Target execution stages for the barrier.
  iree_hal_execution_stage_t target_stage_mask;
  // Barrier flags captured from the original call.
  iree_hal_execution_barrier_flags_t flags;
  // Number of serialized memory barriers.
  iree_host_size_t memory_barrier_count;
  // Serialized memory barriers borrowed from the replay file.
  const iree_hal_replay_memory_barrier_payload_t* memory_barrier_payloads;
  // Number of serialized buffer barriers.
  iree_host_size_t buffer_barrier_count;
  // Serialized buffer barriers borrowed from the replay file.
  const iree_hal_replay_buffer_barrier_payload_t* buffer_barrier_payloads;
} iree_hal_replay_plan_command_buffer_execution_barrier_t;

typedef struct iree_hal_replay_plan_command_buffer_dispatch_t {
  // Captured command buffer object id.
  iree_hal_replay_object_id_t command_buffer_id;
  // Captured executable object id.
  iree_hal_replay_object_id_t executable_id;
  // Captured function ordinal to resolve at replay execution time.
  uint32_t function_ordinal;
  // Dispatch grid and indirect-count configuration.
  iree_hal_dispatch_config_t config;
  // Serialized indirect workgroup count buffer reference.
  iree_hal_replay_buffer_ref_payload_t workgroup_count_ref;
  // Inline constant bytes borrowed from the replay file.
  iree_const_byte_span_t constants;
  // Number of serialized buffer refs in |binding_payloads|.
  iree_host_size_t binding_count;
  // Serialized dispatch bindings borrowed from the replay file.
  const iree_hal_replay_buffer_ref_payload_t* binding_payloads;
  // Dispatch flags captured from the original call.
  iree_hal_dispatch_flags_t flags;
} iree_hal_replay_plan_command_buffer_dispatch_t;

typedef struct iree_hal_replay_plan_queue_execute_t {
  // Captured device object id receiving the queue operation.
  iree_hal_replay_object_id_t device_id;
  // Captured command buffer object id, or NONE for a queue barrier.
  iree_hal_replay_object_id_t command_buffer_id;
  // Queue affinity captured from the original call.
  iree_hal_queue_affinity_t queue_affinity;
  // Queue execution flags captured from the original call.
  iree_hal_execute_flags_t flags;
  // Number of serialized wait semaphore timepoints.
  iree_host_size_t wait_semaphore_count;
  // Serialized wait semaphore timepoints borrowed from the replay file.
  const iree_hal_replay_semaphore_timepoint_payload_t* wait_semaphore_payloads;
  // Number of serialized signal semaphore timepoints.
  iree_host_size_t signal_semaphore_count;
  // Serialized signal semaphore timepoints borrowed from the replay file.
  const iree_hal_replay_semaphore_timepoint_payload_t*
      signal_semaphore_payloads;
  // Number of serialized queue binding entries.
  iree_host_size_t binding_count;
  // Serialized queue binding entries borrowed from the replay file.
  const iree_hal_replay_buffer_ref_payload_t* binding_payloads;
} iree_hal_replay_plan_queue_execute_t;

typedef struct iree_hal_replay_plan_record_t {
  // Parsed replay record borrowed from the original replay file.
  iree_hal_replay_file_record_t file_record;
  // Prepared interpreter opcode for this record.
  iree_hal_replay_plan_record_kind_t kind;
  // Prepared payload for hot replay operations.
  union {
    // Prepared scope marker payload.
    iree_hal_replay_plan_scope_t scope;
    // Prepared command buffer execution barrier payload.
    iree_hal_replay_plan_command_buffer_execution_barrier_t
        command_buffer_execution_barrier;
    // Prepared command buffer dispatch payload.
    iree_hal_replay_plan_command_buffer_dispatch_t command_buffer_dispatch;
    // Prepared queue execute payload.
    iree_hal_replay_plan_queue_execute_t queue_execute;
  } payload;
} iree_hal_replay_plan_record_t;

struct iree_hal_replay_plan_t {
  // Original replay file bytes borrowed by this plan.
  iree_const_byte_span_t file_contents;
  // Host allocator used for the plan allocations.
  iree_allocator_t host_allocator;
  // Dense session-local object table capacity required during execution.
  iree_host_size_t object_capacity;
  // Number of entries in |records|.
  iree_host_size_t record_count;
  // Prepared replay records owned by this plan.
  iree_hal_replay_plan_record_t* records;
};

typedef struct iree_hal_replay_executor_t {
  // Original replay file bytes.
  iree_const_byte_span_t file_contents;
  // Retained topology supplied by the caller.
  iree_hal_device_group_t* device_group;
  // Host allocator used for temporary replay state.
  iree_allocator_t host_allocator;
  // Execution options supplied by the caller.
  const iree_hal_replay_execute_options_t* options;
  // Dense session-local object table indexed by replay object id.
  iree_hal_replay_object_entry_t* objects;
  // Number of entries in |objects|.
  iree_host_size_t object_capacity;
  // Next caller-provided device consumed by a device object record.
  iree_host_size_t next_device_index;
} iree_hal_replay_executor_t;

static void iree_hal_replay_executor_release_entry(
    iree_hal_replay_executor_t* executor,
    iree_hal_replay_object_entry_t* entry) {
  switch (entry->type) {
    case IREE_HAL_REPLAY_OBJECT_TYPE_DEVICE:
      iree_hal_device_release(entry->value.device);
      break;
    case IREE_HAL_REPLAY_OBJECT_TYPE_ALLOCATOR:
      iree_hal_allocator_release(entry->value.allocator);
      break;
    case IREE_HAL_REPLAY_OBJECT_TYPE_BUFFER:
      iree_hal_buffer_release(entry->value.buffer);
      break;
    case IREE_HAL_REPLAY_OBJECT_TYPE_COMMAND_BUFFER:
      iree_hal_command_buffer_release(entry->value.command_buffer);
      break;
    case IREE_HAL_REPLAY_OBJECT_TYPE_EXECUTABLE_CACHE:
      iree_hal_executable_cache_release(entry->value.executable_cache);
      break;
    case IREE_HAL_REPLAY_OBJECT_TYPE_EXECUTABLE:
      iree_hal_executable_release(entry->value.executable.handle);
      iree_allocator_free(executor->host_allocator,
                          entry->value.executable.function_map);
      break;
    case IREE_HAL_REPLAY_OBJECT_TYPE_SEMAPHORE:
      iree_hal_semaphore_release(entry->value.semaphore);
      break;
    case IREE_HAL_REPLAY_OBJECT_TYPE_EVENT:
      iree_hal_event_release(entry->value.event);
      break;
    case IREE_HAL_REPLAY_OBJECT_TYPE_FILE:
      iree_hal_file_release(entry->value.file);
      break;
    default:
      break;
  }
  memset(entry, 0, sizeof(*entry));
}

static void iree_hal_replay_executor_deinitialize(
    iree_hal_replay_executor_t* executor) {
  if (!executor->objects) return;
  for (iree_host_size_t i = executor->object_capacity; i > 0; --i) {
    iree_hal_replay_executor_release_entry(executor, &executor->objects[i - 1]);
  }
  iree_allocator_free(executor->host_allocator, executor->objects);
  executor->objects = NULL;
  executor->object_capacity = 0;
}

static iree_status_t iree_hal_replay_executor_lookup(
    iree_hal_replay_executor_t* executor, iree_hal_replay_object_id_t object_id,
    iree_hal_replay_object_type_t expected_type,
    iree_hal_replay_object_entry_t** out_entry) {
  IREE_ASSERT_ARGUMENT(out_entry);
  *out_entry = NULL;
  if (IREE_UNLIKELY(object_id == IREE_HAL_REPLAY_OBJECT_ID_NONE ||
                    object_id >= executor->object_capacity)) {
    return iree_make_status(IREE_STATUS_NOT_FOUND,
                            "replay object id %" PRIu64 " is not defined",
                            object_id);
  }
  iree_hal_replay_object_entry_t* entry = &executor->objects[object_id];
  if (IREE_UNLIKELY(entry->type != expected_type)) {
    return iree_make_status(
        IREE_STATUS_FAILED_PRECONDITION,
        "replay object id %" PRIu64 " has type %s; expected %s", object_id,
        iree_hal_replay_object_type_string(entry->type),
        iree_hal_replay_object_type_string(expected_type));
  }
  *out_entry = entry;
  return iree_ok_status();
}

static iree_status_t iree_hal_replay_executor_store(
    iree_hal_replay_executor_t* executor, iree_hal_replay_object_id_t object_id,
    iree_hal_replay_object_type_t object_type,
    iree_hal_replay_object_entry_t entry) {
  entry.type = object_type;
  if (IREE_UNLIKELY(object_id == IREE_HAL_REPLAY_OBJECT_ID_NONE ||
                    object_id >= executor->object_capacity)) {
    iree_hal_replay_executor_release_entry(executor, &entry);
    return iree_make_status(IREE_STATUS_OUT_OF_RANGE,
                            "replay object id %" PRIu64
                            " exceeds object table capacity",
                            object_id);
  }
  iree_hal_replay_object_entry_t* existing = &executor->objects[object_id];
  if (existing->type != IREE_HAL_REPLAY_OBJECT_TYPE_NONE) {
    iree_hal_replay_executor_release_entry(executor, &entry);
    return iree_make_status(IREE_STATUS_ALREADY_EXISTS,
                            "replay object id %" PRIu64 " is already assigned",
                            object_id);
  }
  *existing = entry;
  return iree_ok_status();
}

static iree_status_t iree_hal_replay_executor_allocate_function_map(
    iree_hal_replay_executor_t* executor, iree_host_size_t function_count,
    iree_hal_executable_function_t** out_function_map) {
  *out_function_map = NULL;
  if (function_count == 0) return iree_ok_status();
  iree_host_size_t function_map_size = 0;
  if (IREE_UNLIKELY(!iree_host_size_checked_mul(
          function_count, sizeof(iree_hal_executable_function_t),
          &function_map_size))) {
    return iree_make_status(IREE_STATUS_OUT_OF_RANGE,
                            "replay executable function map size overflow");
  }
  return iree_allocator_malloc(executor->host_allocator, function_map_size,
                               (void**)out_function_map);
}

static iree_status_t iree_hal_replay_executor_resolve_function(
    iree_hal_replay_object_id_t executable_id,
    const iree_hal_replay_executable_entry_t* executable,
    uint32_t function_ordinal, iree_hal_executable_function_t* out_function) {
  if (IREE_UNLIKELY(function_ordinal >= executable->function_map_count)) {
    return iree_make_status(
        IREE_STATUS_OUT_OF_RANGE,
        "replay executable %" PRIu64
        " does not define captured function ordinal %" PRIu32,
        executable_id, function_ordinal);
  }
  *out_function = executable->function_map[function_ordinal];
  if (IREE_UNLIKELY(!iree_hal_executable_function_is_valid(*out_function))) {
    return iree_make_status(IREE_STATUS_FAILED_PRECONDITION,
                            "replay executable %" PRIu64
                            " captured function ordinal %" PRIu32
                            " did not resolve",
                            executable_id, function_ordinal);
  }
  return iree_ok_status();
}

static iree_status_t iree_hal_replay_executor_initialize(
    iree_hal_replay_executor_t* executor, const iree_hal_replay_plan_t* plan,
    iree_hal_device_group_t* device_group,
    const iree_hal_replay_execute_options_t* options,
    iree_allocator_t host_allocator) {
  memset(executor, 0, sizeof(*executor));
  executor->file_contents = plan->file_contents;
  executor->device_group = device_group;
  executor->host_allocator = host_allocator;
  executor->options = options;
  executor->object_capacity = plan->object_capacity;
  if (executor->object_capacity == 0) return iree_ok_status();
  iree_host_size_t object_table_size = 0;
  if (IREE_UNLIKELY(!iree_host_size_checked_mul(executor->object_capacity,
                                                sizeof(*executor->objects),
                                                &object_table_size))) {
    return iree_make_status(IREE_STATUS_OUT_OF_RANGE,
                            "replay object table size overflow");
  }
  IREE_RETURN_IF_ERROR(iree_allocator_malloc(host_allocator, object_table_size,
                                             (void**)&executor->objects));
  memset(executor->objects, 0, object_table_size);
  return iree_ok_status();
}

static iree_status_t iree_hal_replay_executor_require_payload(
    const iree_hal_replay_file_record_t* record,
    iree_hal_replay_payload_type_t payload_type,
    iree_host_size_t minimum_length) {
  if (IREE_UNLIKELY(record->header.payload_type != payload_type)) {
    return iree_make_status(
        IREE_STATUS_UNIMPLEMENTED, "replay operation %s requires payload %s",
        iree_hal_replay_operation_code_string(record->header.operation_code),
        iree_hal_replay_payload_type_string(payload_type));
  }
  if (IREE_UNLIKELY(record->payload.data_length < minimum_length)) {
    return iree_make_status(IREE_STATUS_DATA_LOSS,
                            "replay payload %s is too short",
                            iree_hal_replay_payload_type_string(payload_type));
  }
  return iree_ok_status();
}

static iree_status_t iree_hal_replay_executor_scope_event(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_file_record_t* record,
    iree_hal_replay_scope_event_type_t event_type) {
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_require_payload(
      record, IREE_HAL_REPLAY_PAYLOAD_TYPE_REPLAY_SCOPE,
      sizeof(iree_hal_replay_scope_payload_t)));
  iree_hal_replay_scope_payload_t payload;
  memcpy(&payload, record->payload.data, sizeof(payload));
  if (IREE_UNLIKELY(payload.flags != IREE_HAL_REPLAY_SCOPE_FLAG_NONE ||
                    payload.reserved0 != 0 || payload.reserved1 != 0)) {
    return iree_make_status(IREE_STATUS_DATA_LOSS,
                            "replay scope payload reserved fields must be "
                            "zero");
  }
  if (IREE_UNLIKELY(payload.name_length > IREE_HOST_SIZE_MAX ||
                    sizeof(payload) + (iree_host_size_t)payload.name_length !=
                        record->payload.data_length ||
                    payload.name_length == 0)) {
    return iree_make_status(IREE_STATUS_DATA_LOSS,
                            "replay scope payload name length mismatch");
  }

  iree_hal_replay_scope_event_callback_t callback =
      executor->options->scope_event_callback;
  if (!callback.fn) return iree_ok_status();

  iree_hal_replay_scope_event_t event = {
      .sequence_ordinal = record->header.sequence_ordinal,
      .type = event_type,
      .name = iree_make_string_view(
          (const char*)record->payload.data + sizeof(payload),
          (iree_host_size_t)payload.name_length),
  };
  return callback.fn(callback.user_data, &event);
}

static iree_status_t iree_hal_replay_executor_make_buffer_params(
    const iree_hal_replay_allocator_allocate_buffer_payload_t* payload,
    iree_hal_buffer_params_t* out_params) {
  memset(out_params, 0, sizeof(*out_params));
  out_params->usage = payload->usage;
  out_params->type = payload->type;
  out_params->access = payload->access;
  out_params->queue_affinity = payload->queue_affinity;
  out_params->min_alignment = payload->min_alignment;
  return iree_ok_status();
}

static iree_status_t iree_hal_replay_executor_write_buffer_data(
    iree_hal_buffer_t* buffer, iree_device_size_t byte_offset,
    iree_device_size_t byte_length, iree_hal_memory_access_t memory_access,
    iree_const_byte_span_t data) {
  if (data.data_length == 0) {
    return iree_ok_status();
  }
  if (IREE_UNLIKELY(data.data_length > byte_length)) {
    return iree_make_status(IREE_STATUS_DATA_LOSS,
                            "replay buffer data overflows target range");
  }
  iree_hal_buffer_mapping_t mapping;
  IREE_RETURN_IF_ERROR(iree_hal_buffer_map_range(
      buffer, IREE_HAL_MAPPING_MODE_SCOPED, memory_access, byte_offset,
      byte_length, &mapping));
  iree_status_t status = iree_ok_status();
  iree_byte_span_t target_span;
  if (iree_status_is_ok(status)) {
    status =
        iree_hal_buffer_mapping_subspan(&mapping, IREE_HAL_MEMORY_ACCESS_WRITE,
                                        0, data.data_length, &target_span);
  }
  if (iree_status_is_ok(status)) {
    memcpy(target_span.data, data.data, data.data_length);
    status = iree_hal_buffer_mapping_flush_range(&mapping, 0, data.data_length);
  }
  iree_status_t unmap_status = iree_hal_buffer_unmap_range(&mapping);
  return iree_status_join(status, unmap_status);
}

static iree_status_t iree_hal_replay_executor_make_buffer_ref(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_buffer_ref_payload_t* payload,
    iree_hal_buffer_ref_t* out_ref) {
  if (payload->buffer_id == IREE_HAL_REPLAY_OBJECT_ID_NONE) {
    *out_ref = iree_hal_make_indirect_buffer_ref(
        payload->buffer_slot, payload->offset, payload->length);
    return iree_ok_status();
  }
  iree_hal_replay_object_entry_t* entry = NULL;
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_lookup(
      executor, payload->buffer_id, IREE_HAL_REPLAY_OBJECT_TYPE_BUFFER,
      &entry));
  *out_ref = iree_hal_make_buffer_ref(entry->value.buffer, payload->offset,
                                      payload->length);
  return iree_ok_status();
}

typedef struct iree_hal_replay_semaphore_list_storage_t {
  // HAL semaphore list referencing arrays below.
  iree_hal_semaphore_list_t list;
  // Mutable semaphore pointer array used by |list|.
  iree_hal_semaphore_t** semaphores;
  // Mutable semaphore payload value array used by |list|.
  uint64_t* payload_values;
  // Inline storage used for common small semaphore lists.
  struct {
    // Inline semaphore pointer storage.
    iree_hal_semaphore_t*
        semaphores[IREE_HAL_REPLAY_INLINE_SEMAPHORE_LIST_CAPACITY];
    // Inline payload value storage.
    uint64_t payload_values[IREE_HAL_REPLAY_INLINE_SEMAPHORE_LIST_CAPACITY];
  } inline_storage;
  // Heap storage used when a captured list exceeds inline capacity.
  struct {
    // Heap-allocated semaphore pointer storage, or NULL when inline.
    iree_hal_semaphore_t** semaphores;
    // Heap-allocated payload value storage, or NULL when inline.
    uint64_t* payload_values;
  } allocated;
} iree_hal_replay_semaphore_list_storage_t;

static void iree_hal_replay_semaphore_list_storage_deinitialize(
    iree_hal_replay_semaphore_list_storage_t* storage,
    iree_allocator_t host_allocator) {
  iree_allocator_free(host_allocator, storage->allocated.payload_values);
  iree_allocator_free(host_allocator, storage->allocated.semaphores);
  memset(storage, 0, sizeof(*storage));
}

static iree_status_t iree_hal_replay_executor_make_semaphore_list(
    iree_hal_replay_executor_t* executor, iree_const_byte_span_t payloads,
    iree_host_size_t count,
    iree_hal_replay_semaphore_list_storage_t* out_storage) {
  memset(out_storage, 0, sizeof(*out_storage));
  if (count == 0) return iree_ok_status();
  iree_host_size_t expected_length = 0;
  if (IREE_UNLIKELY(!iree_host_size_checked_mul(
                        count,
                        sizeof(iree_hal_replay_semaphore_timepoint_payload_t),
                        &expected_length) ||
                    payloads.data_length != expected_length)) {
    return iree_make_status(IREE_STATUS_DATA_LOSS,
                            "replay semaphore list payload length mismatch");
  }
  iree_status_t status = iree_ok_status();
  if (count <= IREE_HAL_REPLAY_INLINE_SEMAPHORE_LIST_CAPACITY) {
    out_storage->semaphores = out_storage->inline_storage.semaphores;
    out_storage->payload_values = out_storage->inline_storage.payload_values;
  } else {
    status = iree_allocator_malloc(executor->host_allocator,
                                   count * sizeof(*out_storage->semaphores),
                                   (void**)&out_storage->allocated.semaphores);
    if (iree_status_is_ok(status)) {
      status =
          iree_allocator_malloc(executor->host_allocator,
                                count * sizeof(*out_storage->payload_values),
                                (void**)&out_storage->allocated.payload_values);
    }
    if (iree_status_is_ok(status)) {
      out_storage->semaphores = out_storage->allocated.semaphores;
      out_storage->payload_values = out_storage->allocated.payload_values;
    } else {
      iree_hal_replay_semaphore_list_storage_deinitialize(
          out_storage, executor->host_allocator);
      return status;
    }
  }
  const iree_hal_replay_semaphore_timepoint_payload_t* timepoints =
      (const iree_hal_replay_semaphore_timepoint_payload_t*)payloads.data;
  for (iree_host_size_t i = 0; i < count; ++i) {
    iree_hal_replay_object_entry_t* entry = NULL;
    status = iree_hal_replay_executor_lookup(
        executor, timepoints[i].semaphore_id,
        IREE_HAL_REPLAY_OBJECT_TYPE_SEMAPHORE, &entry);
    if (!iree_status_is_ok(status)) break;
    out_storage->semaphores[i] = entry->value.semaphore;
    out_storage->payload_values[i] = timepoints[i].value;
  }
  if (!iree_status_is_ok(status)) {
    iree_hal_replay_semaphore_list_storage_deinitialize(
        out_storage, executor->host_allocator);
    return status;
  }
  out_storage->list.count = count;
  out_storage->list.semaphores = out_storage->semaphores;
  out_storage->list.payload_values = out_storage->payload_values;
  return iree_ok_status();
}

static iree_status_t iree_hal_replay_executor_make_queue_semaphore_lists(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_file_record_t* record, iree_host_size_t header_length,
    uint64_t wait_semaphore_count, uint64_t signal_semaphore_count,
    uint64_t trailing_payload_length,
    iree_hal_replay_semaphore_list_storage_t* out_wait_storage,
    iree_hal_replay_semaphore_list_storage_t* out_signal_storage,
    iree_const_byte_span_t* out_trailing_payload) {
  memset(out_wait_storage, 0, sizeof(*out_wait_storage));
  memset(out_signal_storage, 0, sizeof(*out_signal_storage));
  *out_trailing_payload = iree_make_const_byte_span(NULL, 0);
  iree_host_size_t wait_size = 0;
  iree_host_size_t signal_size = 0;
  iree_host_size_t total_size = 0;
  if (IREE_UNLIKELY(
          wait_semaphore_count > IREE_HOST_SIZE_MAX ||
          signal_semaphore_count > IREE_HOST_SIZE_MAX ||
          trailing_payload_length > IREE_HOST_SIZE_MAX ||
          !iree_host_size_checked_mul(
              (iree_host_size_t)wait_semaphore_count,
              sizeof(iree_hal_replay_semaphore_timepoint_payload_t),
              &wait_size) ||
          !iree_host_size_checked_mul(
              (iree_host_size_t)signal_semaphore_count,
              sizeof(iree_hal_replay_semaphore_timepoint_payload_t),
              &signal_size) ||
          !iree_host_size_checked_add(header_length, wait_size, &total_size) ||
          !iree_host_size_checked_add(total_size, signal_size, &total_size) ||
          !iree_host_size_checked_add(total_size,
                                      (iree_host_size_t)trailing_payload_length,
                                      &total_size) ||
          total_size != record->payload.data_length)) {
    return iree_make_status(IREE_STATUS_DATA_LOSS,
                            "replay queue payload length mismatch");
  }

  const uint8_t* cursor = record->payload.data + header_length;
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_make_semaphore_list(
      executor, iree_make_const_byte_span(cursor, wait_size),
      (iree_host_size_t)wait_semaphore_count, out_wait_storage));
  cursor += wait_size;
  iree_status_t status = iree_hal_replay_executor_make_semaphore_list(
      executor, iree_make_const_byte_span(cursor, signal_size),
      (iree_host_size_t)signal_semaphore_count, out_signal_storage);
  cursor += signal_size;
  if (!iree_status_is_ok(status)) {
    iree_hal_replay_semaphore_list_storage_deinitialize(
        out_wait_storage, executor->host_allocator);
    return status;
  }
  *out_trailing_payload = iree_make_const_byte_span(
      cursor, (iree_host_size_t)trailing_payload_length);
  return iree_ok_status();
}

typedef struct iree_hal_replay_buffer_ref_list_storage_t {
  // HAL buffer ref list referencing |values|.
  iree_hal_buffer_ref_list_t list;
  // Mutable buffer ref array used by |list|.
  iree_hal_buffer_ref_t* values;
  // Inline storage used for common small binding lists.
  struct {
    // Inline buffer ref storage.
    iree_hal_buffer_ref_t
        values[IREE_HAL_REPLAY_INLINE_BUFFER_REF_LIST_CAPACITY];
  } inline_storage;
  // Heap storage used when a captured list exceeds inline capacity.
  struct {
    // Heap-allocated buffer ref storage, or NULL when inline.
    iree_hal_buffer_ref_t* values;
  } allocated;
} iree_hal_replay_buffer_ref_list_storage_t;

static void iree_hal_replay_buffer_ref_list_storage_deinitialize(
    iree_hal_replay_buffer_ref_list_storage_t* storage,
    iree_allocator_t host_allocator) {
  iree_allocator_free(host_allocator, storage->allocated.values);
  memset(storage, 0, sizeof(*storage));
}

static iree_status_t iree_hal_replay_buffer_ref_list_storage_initialize(
    iree_hal_replay_executor_t* executor, iree_host_size_t count,
    iree_hal_replay_buffer_ref_list_storage_t* out_storage) {
  memset(out_storage, 0, sizeof(*out_storage));
  if (count == 0) return iree_ok_status();
  if (count <= IREE_HAL_REPLAY_INLINE_BUFFER_REF_LIST_CAPACITY) {
    out_storage->values = out_storage->inline_storage.values;
  } else {
    iree_host_size_t values_size = 0;
    if (IREE_UNLIKELY(!iree_host_size_checked_mul(
            count, sizeof(*out_storage->values), &values_size))) {
      return iree_make_status(IREE_STATUS_OUT_OF_RANGE,
                              "replay buffer ref list count overflow");
    }
    IREE_RETURN_IF_ERROR(
        iree_allocator_malloc(executor->host_allocator, values_size,
                              (void**)&out_storage->allocated.values));
    out_storage->values = out_storage->allocated.values;
  }
  out_storage->list.count = count;
  out_storage->list.values = out_storage->values;
  return iree_ok_status();
}

typedef struct iree_hal_replay_buffer_binding_table_storage_t {
  // HAL buffer binding table referencing |bindings|.
  iree_hal_buffer_binding_table_t table;
  // Mutable buffer binding array used by |table|.
  iree_hal_buffer_binding_t* bindings;
  // Inline storage used for common small binding tables.
  struct {
    // Inline buffer binding storage.
    iree_hal_buffer_binding_t
        bindings[IREE_HAL_REPLAY_INLINE_BUFFER_BINDING_TABLE_CAPACITY];
  } inline_storage;
  // Heap storage used when a captured table exceeds inline capacity.
  struct {
    // Heap-allocated buffer binding storage, or NULL when inline.
    iree_hal_buffer_binding_t* bindings;
  } allocated;
} iree_hal_replay_buffer_binding_table_storage_t;

static void iree_hal_replay_buffer_binding_table_storage_deinitialize(
    iree_hal_replay_buffer_binding_table_storage_t* storage,
    iree_allocator_t host_allocator) {
  iree_allocator_free(host_allocator, storage->allocated.bindings);
  memset(storage, 0, sizeof(*storage));
}

static iree_status_t iree_hal_replay_buffer_binding_table_storage_initialize(
    iree_hal_replay_executor_t* executor, iree_host_size_t count,
    iree_hal_replay_buffer_binding_table_storage_t* out_storage) {
  memset(out_storage, 0, sizeof(*out_storage));
  if (count == 0) return iree_ok_status();
  if (count <= IREE_HAL_REPLAY_INLINE_BUFFER_BINDING_TABLE_CAPACITY) {
    out_storage->bindings = out_storage->inline_storage.bindings;
  } else {
    iree_host_size_t bindings_size = 0;
    if (IREE_UNLIKELY(!iree_host_size_checked_mul(
            count, sizeof(*out_storage->bindings), &bindings_size))) {
      return iree_make_status(IREE_STATUS_OUT_OF_RANGE,
                              "replay buffer binding table count overflow");
    }
    IREE_RETURN_IF_ERROR(
        iree_allocator_malloc(executor->host_allocator, bindings_size,
                              (void**)&out_storage->allocated.bindings));
    out_storage->bindings = out_storage->allocated.bindings;
  }
  out_storage->table.count = count;
  out_storage->table.bindings = out_storage->bindings;
  return iree_ok_status();
}

static iree_status_t iree_hal_replay_executor_flush_and_wait(
    iree_hal_device_t* device, iree_hal_queue_affinity_t queue_affinity,
    const iree_hal_semaphore_list_t signal_list) {
  IREE_RETURN_IF_ERROR(iree_hal_device_queue_flush(device, queue_affinity));
  if (signal_list.count == 0) return iree_ok_status();
  return iree_hal_semaphore_list_wait(signal_list, iree_infinite_timeout(),
                                      IREE_ASYNC_WAIT_FLAG_NONE);
}

static iree_status_t iree_hal_replay_executor_store_device(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_file_record_t* record) {
  if (executor->next_device_index >=
      iree_hal_device_group_device_count(executor->device_group)) {
    return iree_make_status(IREE_STATUS_OUT_OF_RANGE,
                            "replay references more devices than provided");
  }
  iree_hal_device_t* device = iree_hal_device_group_device_at(
      executor->device_group, executor->next_device_index++);
  iree_hal_device_retain(device);
  iree_hal_replay_object_entry_t entry = {.value.device = device};
  return iree_hal_replay_executor_store(executor, record->header.object_id,
                                        IREE_HAL_REPLAY_OBJECT_TYPE_DEVICE,
                                        entry);
}

static iree_status_t iree_hal_replay_executor_store_allocator(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_file_record_t* record) {
  iree_hal_replay_object_entry_t* device_entry = NULL;
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_lookup(
      executor, record->header.device_id, IREE_HAL_REPLAY_OBJECT_TYPE_DEVICE,
      &device_entry));
  iree_hal_allocator_t* allocator =
      iree_hal_device_allocator(device_entry->value.device);
  iree_hal_allocator_retain(allocator);
  iree_hal_replay_object_entry_t entry = {.value.allocator = allocator};
  return iree_hal_replay_executor_store(executor, record->header.object_id,
                                        IREE_HAL_REPLAY_OBJECT_TYPE_ALLOCATOR,
                                        entry);
}

static iree_status_t iree_hal_replay_executor_replay_object(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_file_record_t* record) {
  switch (record->header.object_type) {
    case IREE_HAL_REPLAY_OBJECT_TYPE_DEVICE:
      return iree_hal_replay_executor_store_device(executor, record);
    case IREE_HAL_REPLAY_OBJECT_TYPE_ALLOCATOR:
      return iree_hal_replay_executor_store_allocator(executor, record);
    default:
      return iree_ok_status();
  }
}

static iree_status_t iree_hal_replay_executor_make_function_map_from_metadata(
    iree_hal_replay_executor_t* executor,
    iree_hal_replay_object_id_t executable_id,
    iree_const_byte_span_t executable_metadata,
    iree_hal_executable_t* executable, iree_host_size_t* out_function_map_count,
    iree_hal_executable_function_t** out_function_map) {
  *out_function_map_count = 0;
  *out_function_map = NULL;

  if (IREE_UNLIKELY(executable_metadata.data_length <
                    sizeof(iree_hal_replay_executable_metadata_header_t))) {
    return iree_make_status(IREE_STATUS_DATA_LOSS,
                            "replay executable metadata is too short");
  }
  iree_hal_replay_executable_metadata_header_t header;
  memcpy(&header, executable_metadata.data, sizeof(header));
  if (IREE_UNLIKELY(header.reserved1 != 0)) {
    return iree_make_status(IREE_STATUS_DATA_LOSS,
                            "replay executable metadata reserved fields must "
                            "be zero");
  }
  if (IREE_UNLIKELY(header.function_count > IREE_HOST_SIZE_MAX ||
                    header.function_count > UINT32_MAX ||
                    header.parameter_count > IREE_HOST_SIZE_MAX ||
                    header.function_name_storage_length > IREE_HOST_SIZE_MAX)) {
    return iree_make_status(IREE_STATUS_OUT_OF_RANGE,
                            "replay executable metadata count overflow");
  }
  const iree_host_size_t captured_function_count =
      (iree_host_size_t)header.function_count;
  const iree_host_size_t captured_parameter_count =
      (iree_host_size_t)header.parameter_count;
  const iree_host_size_t captured_name_storage_length =
      (iree_host_size_t)header.function_name_storage_length;

  iree_host_size_t function_metadata_size = 0;
  iree_host_size_t parameter_metadata_size = 0;
  iree_host_size_t expected_length = 0;
  if (IREE_UNLIKELY(!iree_host_size_checked_mul(
                        captured_function_count,
                        sizeof(iree_hal_replay_executable_function_metadata_t),
                        &function_metadata_size) ||
                    !iree_host_size_checked_mul(
                        captured_parameter_count,
                        sizeof(iree_hal_replay_executable_parameter_metadata_t),
                        &parameter_metadata_size) ||
                    !iree_host_size_checked_add(
                        sizeof(iree_hal_replay_executable_metadata_header_t),
                        function_metadata_size, &expected_length) ||
                    !iree_host_size_checked_add(expected_length,
                                                parameter_metadata_size,
                                                &expected_length) ||
                    !iree_host_size_checked_add(expected_length,
                                                captured_name_storage_length,
                                                &expected_length) ||
                    expected_length != executable_metadata.data_length)) {
    return iree_make_status(IREE_STATUS_DATA_LOSS,
                            "replay executable metadata length mismatch");
  }

  const uint8_t* function_metadata_data =
      executable_metadata.data +
      sizeof(iree_hal_replay_executable_metadata_header_t);
  const uint8_t* parameter_metadata_data =
      function_metadata_data + function_metadata_size;
  const iree_hal_replay_executable_function_metadata_t* function_metadata =
      (const iree_hal_replay_executable_function_metadata_t*)
          function_metadata_data;
  const iree_hal_replay_executable_parameter_metadata_t* parameter_metadata =
      (const iree_hal_replay_executable_parameter_metadata_t*)
          parameter_metadata_data;
  const char* name_storage =
      (const char*)(parameter_metadata_data + parameter_metadata_size);

  iree_hal_executable_function_t* function_map = NULL;
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_allocate_function_map(
      executor, captured_function_count, &function_map));

  iree_host_size_t parameter_index = 0;
  iree_host_size_t name_offset = 0;
  iree_status_t status = iree_ok_status();
  for (iree_host_size_t i = 0;
       i < captured_function_count && iree_status_is_ok(status); ++i) {
    const iree_hal_replay_executable_function_metadata_t* captured =
        &function_metadata[i];
    if (IREE_UNLIKELY(captured->parameter_count >
                          captured_parameter_count - parameter_index ||
                      captured->name_length >
                          captured_name_storage_length - name_offset)) {
      status = iree_make_status(IREE_STATUS_DATA_LOSS,
                                "replay executable metadata count mismatch");
    }

    iree_string_view_t captured_name = iree_string_view_empty();
    if (iree_status_is_ok(status)) {
      captured_name = iree_make_string_view(
          name_storage + name_offset, (iree_host_size_t)captured->name_length);
      name_offset += captured->name_length;
    }

    iree_hal_executable_function_t function =
        iree_hal_executable_function_invalid();
    if (iree_status_is_ok(status)) {
      if (iree_string_view_is_empty(captured_name)) {
        status =
            iree_make_status(IREE_STATUS_DATA_LOSS,
                             "replay executable %" PRIu64
                             " captured function %" PRIhsz " is missing a name",
                             executable_id, i);
      }
    }
    if (iree_status_is_ok(status)) {
      status = iree_hal_executable_lookup_function_by_name(
          executable, captured_name, &function);
      if (!iree_status_is_ok(status)) {
        status = iree_status_annotate_f(
            status,
            "looking up replay function '%.*s' for executable %" PRIu64
            " captured function %" PRIhsz,
            (int)captured_name.size, captured_name.data, executable_id, i);
      }
    }
    for (iree_host_size_t j = 0; j < i && iree_status_is_ok(status); ++j) {
      if (IREE_UNLIKELY(function_map[j].value == function.value)) {
        status = iree_make_status(
            IREE_STATUS_FAILED_PRECONDITION,
            "replay executable %" PRIu64
            " captured function '%.*s' resolves to the same loaded function "
            "token as captured function %" PRIhsz,
            executable_id, (int)captured_name.size, captured_name.data, j);
      }
    }

    iree_hal_executable_function_info_t loaded_info;
    if (iree_status_is_ok(status)) {
      status =
          iree_hal_executable_function_info(executable, function, &loaded_info);
    }
    if (iree_status_is_ok(status) &&
        (captured->flags != loaded_info.flags ||
         captured->constant_count != loaded_info.constant_count ||
         captured->binding_count != loaded_info.binding_count ||
         captured->parameter_count != loaded_info.parameter_count ||
         captured->workgroup_size[0] != loaded_info.workgroup_size[0] ||
         captured->workgroup_size[1] != loaded_info.workgroup_size[1] ||
         captured->workgroup_size[2] != loaded_info.workgroup_size[2])) {
      status = iree_make_status(
          IREE_STATUS_FAILED_PRECONDITION,
          "replay executable %" PRIu64 " function %" PRIhsz
          " ABI mismatch: captured=(flags=0x%016" PRIx64
          " constants=%u bindings=%u parameters=%u workgroup_size=[%u,%u,%u]) "
          "loaded=(flags=0x%016" PRIx64
          " constants=%u bindings=%u parameters=%u workgroup_size=[%u,%u,%u])",
          executable_id, i, captured->flags, (uint32_t)captured->constant_count,
          (uint32_t)captured->binding_count,
          (uint32_t)captured->parameter_count, captured->workgroup_size[0],
          captured->workgroup_size[1], captured->workgroup_size[2],
          loaded_info.flags, (uint32_t)loaded_info.constant_count,
          (uint32_t)loaded_info.binding_count,
          (uint32_t)loaded_info.parameter_count, loaded_info.workgroup_size[0],
          loaded_info.workgroup_size[1], loaded_info.workgroup_size[2]);
    }

    iree_hal_executable_function_parameter_t* loaded_parameters = NULL;
    if (iree_status_is_ok(status) && captured->parameter_count != 0) {
      iree_host_size_t parameter_size = 0;
      if (IREE_UNLIKELY(!iree_host_size_checked_mul(
              captured->parameter_count,
              sizeof(iree_hal_executable_function_parameter_t),
              &parameter_size))) {
        status = iree_make_status(IREE_STATUS_OUT_OF_RANGE,
                                  "executable parameter metadata is too large");
      }
      if (iree_status_is_ok(status)) {
        status = iree_allocator_malloc(executor->host_allocator, parameter_size,
                                       (void**)&loaded_parameters);
      }
      if (iree_status_is_ok(status)) {
        status = iree_hal_executable_function_parameters(
            executable, function, captured->parameter_count, loaded_parameters);
      }
    }
    for (iree_host_size_t j = 0;
         j < captured->parameter_count && iree_status_is_ok(status); ++j) {
      const iree_hal_replay_executable_parameter_metadata_t*
          captured_parameter = &parameter_metadata[parameter_index + j];
      if (IREE_UNLIKELY(captured_parameter->reserved0 != 0)) {
        status = iree_make_status(IREE_STATUS_DATA_LOSS,
                                  "replay executable parameter metadata "
                                  "reserved fields must be zero");
        break;
      }
      const iree_hal_executable_function_parameter_t* loaded_parameter =
          &loaded_parameters[j];
      if (captured_parameter->type != loaded_parameter->type ||
          captured_parameter->size != loaded_parameter->size ||
          captured_parameter->flags != loaded_parameter->flags ||
          captured_parameter->offset != loaded_parameter->offset) {
        status = iree_make_status(
            IREE_STATUS_FAILED_PRECONDITION,
            "replay executable %" PRIu64 " function %" PRIhsz
            " parameter %" PRIhsz
            " ABI mismatch: captured=(type=%u size=%u flags=0x%04x "
            "offset=%u) loaded=(type=%u size=%u flags=0x%04x offset=%u)",
            executable_id, i, j, (uint32_t)captured_parameter->type,
            (uint32_t)captured_parameter->size,
            (uint32_t)captured_parameter->flags,
            (uint32_t)captured_parameter->offset,
            (uint32_t)loaded_parameter->type, (uint32_t)loaded_parameter->size,
            (uint32_t)loaded_parameter->flags,
            (uint32_t)loaded_parameter->offset);
      }
    }
    parameter_index += captured->parameter_count;
    iree_allocator_free(executor->host_allocator, loaded_parameters);
    if (iree_status_is_ok(status)) {
      function_map[i] = function;
    }
  }
  if (iree_status_is_ok(status) &&
      IREE_UNLIKELY(parameter_index != captured_parameter_count ||
                    name_offset != captured_name_storage_length)) {
    status = iree_make_status(IREE_STATUS_DATA_LOSS,
                              "replay executable metadata count mismatch");
  }
  if (iree_status_is_ok(status)) {
    *out_function_map_count = captured_function_count;
    *out_function_map = function_map;
  } else {
    iree_allocator_free(executor->host_allocator, function_map);
  }
  return status;
}

static iree_status_t iree_hal_replay_executor_prepare_captured_executable(
    iree_hal_executable_cache_t* executable_cache,
    const iree_hal_executable_params_t* captured_params,
    iree_hal_executable_t** out_executable) {
  return iree_hal_executable_cache_prepare_executable(
      executable_cache, captured_params, out_executable);
}

static iree_status_t iree_hal_replay_executor_prepare_substitute_executable(
    iree_hal_executable_cache_t* executable_cache,
    const iree_hal_executable_params_t* captured_params,
    const iree_hal_replay_executable_substitution_t* substitution,
    iree_hal_executable_t** out_executable) {
  IREE_ASSERT_ARGUMENT(out_executable);
  *out_executable = NULL;
  if (IREE_UNLIKELY(substitution->executable_data.data_length != 0 &&
                    !substitution->executable_data.data)) {
    return iree_make_status(
        IREE_STATUS_INVALID_ARGUMENT,
        "substituted executable data has non-zero length but no data pointer");
  }
  if (IREE_UNLIKELY(substitution->executable_data.data_length == 0)) {
    return iree_make_status(IREE_STATUS_INVALID_ARGUMENT,
                            "substituted executable data is empty");
  }

  iree_hal_executable_params_t replacement_params = *captured_params;
  replacement_params.caching_mode &=
      ~IREE_HAL_EXECUTABLE_CACHING_MODE_ALIAS_PROVIDED_DATA;
  replacement_params.executable_data = substitution->executable_data;

  char inferred_format[IREE_HAL_REPLAY_EXECUTABLE_FORMAT_CAPACITY] = {0};
  if (iree_string_view_is_empty(substitution->executable_format)) {
    iree_host_size_t inferred_size = 0;
    IREE_RETURN_IF_ERROR(iree_status_annotate(
        iree_hal_executable_cache_infer_format(
            executable_cache, replacement_params.caching_mode,
            replacement_params.executable_data, sizeof(inferred_format),
            inferred_format, &inferred_size),
        iree_make_cstring_view(
            "inferring substituted executable format; provide an explicit "
            "format if the target cache cannot infer one")));
    replacement_params.executable_format =
        iree_make_cstring_view(inferred_format);
  } else {
    replacement_params.executable_format = substitution->executable_format;
  }
  return iree_hal_executable_cache_prepare_executable(
      executable_cache, &replacement_params, out_executable);
}

static iree_status_t iree_hal_replay_executor_prepare_executable(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_file_record_t* record) {
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_require_payload(
      record, IREE_HAL_REPLAY_PAYLOAD_TYPE_EXECUTABLE_PREPARE,
      sizeof(iree_hal_replay_executable_prepare_payload_t)));
  iree_hal_replay_executable_prepare_payload_t payload;
  memcpy(&payload, record->payload.data, sizeof(payload));

  iree_host_size_t constant_bytes = 0;
  iree_host_size_t data_offset = 0;
  iree_host_size_t constants_offset = 0;
  iree_host_size_t metadata_offset = 0;
  iree_host_size_t expected_length = 0;
  if (IREE_UNLIKELY(
          payload.executable_data_length > IREE_HOST_SIZE_MAX ||
          payload.constant_count > IREE_HOST_SIZE_MAX ||
          !iree_host_size_checked_mul((iree_host_size_t)payload.constant_count,
                                      sizeof(uint32_t), &constant_bytes) ||
          !iree_host_size_checked_add(sizeof(payload),
                                      payload.executable_format_length,
                                      &data_offset) ||
          !iree_host_size_checked_add(
              data_offset, (iree_host_size_t)payload.executable_data_length,
              &constants_offset) ||
          !iree_host_size_checked_add(constants_offset, constant_bytes,
                                      &metadata_offset) ||
          !iree_host_size_checked_add(
              metadata_offset,
              (iree_host_size_t)payload.executable_metadata_length,
              &expected_length) ||
          expected_length != record->payload.data_length)) {
    return iree_make_status(
        IREE_STATUS_DATA_LOSS,
        "replay executable prepare payload length mismatch");
  }

  iree_hal_replay_object_entry_t* cache_entry = NULL;
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_lookup(
      executor, record->header.object_id,
      IREE_HAL_REPLAY_OBJECT_TYPE_EXECUTABLE_CACHE, &cache_entry));
  iree_hal_executable_params_t params;
  memset(&params, 0, sizeof(params));
  params.queue_affinity = payload.queue_affinity;
  params.executable_format =
      iree_make_string_view((const char*)record->payload.data + sizeof(payload),
                            payload.executable_format_length);
  params.executable_data = iree_make_const_byte_span(
      record->payload.data + data_offset,
      (iree_host_size_t)payload.executable_data_length);
  params.constant_count = (iree_host_size_t)payload.constant_count;
  params.constants = (const uint32_t*)(record->payload.data + constants_offset);
  params.caching_mode = payload.caching_mode &
                        ~IREE_HAL_EXECUTABLE_CACHING_MODE_ALIAS_PROVIDED_DATA;
  iree_const_byte_span_t executable_metadata = iree_make_const_byte_span(
      record->payload.data + metadata_offset,
      (iree_host_size_t)payload.executable_metadata_length);
  if (IREE_UNLIKELY(executable_metadata.data_length == 0)) {
    return iree_make_status(
        IREE_STATUS_DATA_LOSS,
        "replay executable prepare records require function metadata");
  }

  uint8_t* executable_data_storage = NULL;
  uint32_t* constants_storage = NULL;
  iree_hal_executable_t* executable = NULL;
  iree_host_size_t function_map_count = 0;
  iree_hal_executable_function_t* function_map = NULL;
  bool function_map_ready = false;
  iree_status_t status = iree_ok_status();
  if (payload.executable_data_length != 0) {
    status = iree_allocator_malloc(executor->host_allocator,
                                   params.executable_data.data_length,
                                   (void**)&executable_data_storage);
    if (iree_status_is_ok(status)) {
      memcpy(executable_data_storage, params.executable_data.data,
             params.executable_data.data_length);
      params.executable_data = iree_make_const_byte_span(
          executable_data_storage, params.executable_data.data_length);
    }
  }
  if (iree_status_is_ok(status) && constant_bytes != 0) {
    status = iree_allocator_malloc(executor->host_allocator, constant_bytes,
                                   (void**)&constants_storage);
    if (iree_status_is_ok(status)) {
      memcpy(constants_storage, record->payload.data + constants_offset,
             constant_bytes);
      params.constants = constants_storage;
    }
  }
  if (iree_status_is_ok(status) &&
      executor->options->executable_substitution_callback.fn) {
    iree_hal_replay_executable_substitution_request_t request = {
        .sequence_ordinal = record->header.sequence_ordinal,
        .device_id = record->header.device_id,
        .executable_cache_id = record->header.object_id,
        .executable_id = record->header.related_object_id,
        .captured_params = &params,
    };
    iree_hal_replay_executable_substitution_t substitution;
    memset(&substitution, 0, sizeof(substitution));
    status = executor->options->executable_substitution_callback.fn(
        executor->options->executable_substitution_callback.user_data, &request,
        &substitution);
    if (iree_status_is_ok(status) && substitution.substitute) {
      status = iree_status_annotate_f(
          iree_hal_replay_executor_prepare_substitute_executable(
              cache_entry->value.executable_cache, &params, &substitution,
              &executable),
          "preparing substitute for captured executable %" PRIu64
          " from '%.*s'",
          record->header.related_object_id, (int)substitution.source.size,
          substitution.source.data);
      if (iree_status_is_ok(status)) {
        status = iree_hal_replay_executor_make_function_map_from_metadata(
            executor, record->header.related_object_id, executable_metadata,
            executable, &function_map_count, &function_map);
        function_map_ready = iree_status_is_ok(status);
      }
      if (!iree_status_is_ok(status)) {
        iree_hal_executable_release(executable);
        executable = NULL;
      }
    }
  }
  if (iree_status_is_ok(status) && !executable) {
    status = iree_hal_replay_executor_prepare_captured_executable(
        cache_entry->value.executable_cache, &params, &executable);
  }
  if (iree_status_is_ok(status) && !function_map_ready) {
    status = iree_hal_replay_executor_make_function_map_from_metadata(
        executor, record->header.related_object_id, executable_metadata,
        executable, &function_map_count, &function_map);
    function_map_ready = iree_status_is_ok(status);
  }
  if (iree_status_is_ok(status)) {
    iree_hal_replay_object_entry_t entry = {
        .value.executable = {
            .handle = executable,
            .function_map_count = function_map_count,
            .function_map = function_map,
        }};
    status = iree_hal_replay_executor_store(
        executor, record->header.related_object_id,
        IREE_HAL_REPLAY_OBJECT_TYPE_EXECUTABLE, entry);
    executable = NULL;
    function_map = NULL;
    function_map_count = 0;
  } else {
    iree_hal_executable_release(executable);
    iree_allocator_free(executor->host_allocator, function_map);
  }
  iree_allocator_free(executor->host_allocator, constants_storage);
  iree_allocator_free(executor->host_allocator, executable_data_storage);
  return status;
}

static iree_status_t iree_hal_replay_executor_create_executable_cache(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_file_record_t* record) {
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_require_payload(
      record, IREE_HAL_REPLAY_PAYLOAD_TYPE_EXECUTABLE_CACHE_OBJECT,
      sizeof(iree_hal_replay_executable_cache_object_payload_t)));
  iree_hal_replay_executable_cache_object_payload_t payload;
  memcpy(&payload, record->payload.data, sizeof(payload));
  if (IREE_UNLIKELY(payload.identifier_length >
                    record->payload.data_length - sizeof(payload))) {
    return iree_make_status(IREE_STATUS_DATA_LOSS,
                            "replay executable cache identifier overflow");
  }
  iree_hal_replay_object_entry_t* device_entry = NULL;
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_lookup(
      executor, record->header.object_id, IREE_HAL_REPLAY_OBJECT_TYPE_DEVICE,
      &device_entry));
  iree_string_view_t identifier =
      iree_make_string_view((const char*)record->payload.data + sizeof(payload),
                            (iree_host_size_t)payload.identifier_length);
  iree_hal_executable_cache_t* executable_cache = NULL;
  IREE_RETURN_IF_ERROR(iree_hal_executable_cache_create(
      device_entry->value.device, identifier, &executable_cache));
  iree_hal_replay_object_entry_t entry = {.value.executable_cache =
                                              executable_cache};
  return iree_hal_replay_executor_store(
      executor, record->header.related_object_id,
      IREE_HAL_REPLAY_OBJECT_TYPE_EXECUTABLE_CACHE, entry);
}

static iree_status_t iree_hal_replay_executor_create_command_buffer(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_file_record_t* record) {
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_require_payload(
      record, IREE_HAL_REPLAY_PAYLOAD_TYPE_COMMAND_BUFFER_OBJECT,
      sizeof(iree_hal_replay_command_buffer_object_payload_t)));
  iree_hal_replay_command_buffer_object_payload_t payload;
  memcpy(&payload, record->payload.data, sizeof(payload));
  iree_hal_replay_object_entry_t* device_entry = NULL;
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_lookup(
      executor, record->header.object_id, IREE_HAL_REPLAY_OBJECT_TYPE_DEVICE,
      &device_entry));
  iree_hal_command_buffer_t* command_buffer = NULL;
  IREE_RETURN_IF_ERROR(iree_hal_command_buffer_create(
      device_entry->value.device, payload.mode, payload.command_categories,
      payload.queue_affinity, (iree_host_size_t)payload.binding_capacity,
      &command_buffer));
  iree_hal_replay_object_entry_t entry = {.value.command_buffer =
                                              command_buffer};
  return iree_hal_replay_executor_store(
      executor, record->header.related_object_id,
      IREE_HAL_REPLAY_OBJECT_TYPE_COMMAND_BUFFER, entry);
}

static iree_status_t iree_hal_replay_executor_create_semaphore(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_file_record_t* record) {
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_require_payload(
      record, IREE_HAL_REPLAY_PAYLOAD_TYPE_SEMAPHORE_OBJECT,
      sizeof(iree_hal_replay_semaphore_object_payload_t)));
  iree_hal_replay_semaphore_object_payload_t payload;
  memcpy(&payload, record->payload.data, sizeof(payload));
  iree_hal_replay_object_entry_t* device_entry = NULL;
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_lookup(
      executor, record->header.object_id, IREE_HAL_REPLAY_OBJECT_TYPE_DEVICE,
      &device_entry));
  iree_hal_semaphore_t* semaphore = NULL;
  IREE_RETURN_IF_ERROR(iree_hal_semaphore_create(
      device_entry->value.device, payload.queue_affinity, payload.initial_value,
      payload.flags, &semaphore));
  iree_hal_replay_object_entry_t entry = {.value.semaphore = semaphore};
  return iree_hal_replay_executor_store(
      executor, record->header.related_object_id,
      IREE_HAL_REPLAY_OBJECT_TYPE_SEMAPHORE, entry);
}

static iree_status_t iree_hal_replay_executor_create_event(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_file_record_t* record) {
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_require_payload(
      record, IREE_HAL_REPLAY_PAYLOAD_TYPE_EVENT_OBJECT,
      sizeof(iree_hal_replay_event_object_payload_t)));
  iree_hal_replay_event_object_payload_t payload;
  memcpy(&payload, record->payload.data, sizeof(payload));
  iree_hal_replay_object_entry_t* device_entry = NULL;
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_lookup(
      executor, record->header.object_id, IREE_HAL_REPLAY_OBJECT_TYPE_DEVICE,
      &device_entry));
  iree_hal_event_t* event = NULL;
  IREE_RETURN_IF_ERROR(iree_hal_event_create(device_entry->value.device,
                                             payload.queue_affinity,
                                             payload.flags, &event));
  iree_hal_replay_object_entry_t entry = {.value.event = event};
  return iree_hal_replay_executor_store(
      executor, record->header.related_object_id,
      IREE_HAL_REPLAY_OBJECT_TYPE_EVENT, entry);
}

static iree_status_t iree_hal_replay_executor_validate_file_reference(
    const iree_hal_replay_file_object_payload_t* payload,
    iree_string_view_t captured_path, iree_string_view_t resolved_path,
    iree_io_file_handle_t* handle, iree_hal_file_t* file) {
  const uint64_t file_length = iree_hal_file_length(file);
  if (IREE_UNLIKELY(payload->file_length != file_length)) {
    return iree_make_status(
        IREE_STATUS_FAILED_PRECONDITION,
        "external replay file length mismatch for '%.*s' captured as '%.*s': "
        "captured=%" PRIu64 " current=%" PRIu64
        "; restore the matching file or fix --replay_file_remap",
        (int)resolved_path.size, resolved_path.data, (int)captured_path.size,
        captured_path.data, payload->file_length, file_length);
  }

  switch (payload->validation_type) {
    case IREE_HAL_REPLAY_FILE_VALIDATION_TYPE_NONE:
      return iree_ok_status();
    case IREE_HAL_REPLAY_FILE_VALIDATION_TYPE_IDENTITY:
      break;
    case IREE_HAL_REPLAY_FILE_VALIDATION_TYPE_CONTENT_DIGEST:
      break;
    default:
      return iree_make_status(IREE_STATUS_UNIMPLEMENTED,
                              "external replay file validation type %" PRIu32
                              " is not executable",
                              payload->validation_type);
  }

#if IREE_FILE_IO_ENABLE && \
    (defined(IREE_PLATFORM_ANDROID) || defined(IREE_PLATFORM_LINUX))
  iree_io_file_handle_primitive_t primitive =
      iree_io_file_handle_primitive(handle);
  if (IREE_UNLIKELY(primitive.type != IREE_IO_FILE_HANDLE_TYPE_FD)) {
    return iree_make_status(
        IREE_STATUS_FAILED_PRECONDITION,
        "external replay file validation requires an fd-backed file");
  }
  if (payload->validation_type ==
      IREE_HAL_REPLAY_FILE_VALIDATION_TYPE_CONTENT_DIGEST) {
    if (IREE_UNLIKELY(payload->digest_type !=
                      IREE_HAL_REPLAY_DIGEST_TYPE_FNV1A_64)) {
      return iree_make_status(IREE_STATUS_UNIMPLEMENTED,
                              "external replay file digest type %" PRIu32
                              " is not executable",
                              (uint32_t)payload->digest_type);
    }
    uint64_t state = iree_hal_replay_digest_fnv1a64_initialize();
    uint64_t offset = 0;
    uint8_t buffer[64 * 1024];
    while (offset < payload->file_length) {
      uint64_t chunk_length = payload->file_length - offset;
      if (chunk_length > sizeof(buffer)) chunk_length = sizeof(buffer);
      ssize_t read_length = pread(primitive.value.fd, buffer,
                                  (size_t)chunk_length, (off_t)offset);
      if (read_length < 0 && errno == EINTR) continue;
      if (read_length <= 0) {
        return iree_make_status(
            IREE_STATUS_UNAVAILABLE,
            "unable to read external replay file '%.*s' for digest validation",
            (int)resolved_path.size, resolved_path.data);
      }
      state = iree_hal_replay_digest_fnv1a64_update(
          state,
          iree_make_const_byte_span(buffer, (iree_host_size_t)read_length));
      offset += (uint64_t)read_length;
    }
    const uint64_t expected_digest =
        iree_hal_replay_digest_load_fnv1a64(payload->digest);
    if (IREE_UNLIKELY(state != expected_digest)) {
      return iree_make_status(
          IREE_STATUS_FAILED_PRECONDITION,
          "external replay file digest mismatch for '%.*s' captured as "
          "'%.*s': expected=0x%016" PRIx64 " actual=0x%016" PRIx64
          "; restore the matching file or fix --replay_file_remap",
          (int)resolved_path.size, resolved_path.data, (int)captured_path.size,
          captured_path.data, expected_digest, state);
    }
    return iree_ok_status();
  }

  if (payload->file_device == 0 && payload->file_inode == 0 &&
      payload->file_mtime_ns == 0) {
    return iree_ok_status();
  }
  struct stat file_stat;
  if (IREE_UNLIKELY(fstat(primitive.value.fd, &file_stat) != 0)) {
    return iree_make_status(IREE_STATUS_UNAVAILABLE,
                            "unable to stat external replay file '%.*s'",
                            (int)resolved_path.size, resolved_path.data);
  }
  const uint64_t file_device = (uint64_t)file_stat.st_dev;
  const uint64_t file_inode = (uint64_t)file_stat.st_ino;
  const uint64_t file_mtime_ns =
      ((uint64_t)file_stat.st_mtim.tv_sec * 1000000000ull) +
      (uint64_t)file_stat.st_mtim.tv_nsec;
  if (IREE_UNLIKELY(payload->file_device != file_device ||
                    payload->file_inode != file_inode ||
                    payload->file_mtime_ns != file_mtime_ns)) {
    return iree_make_status(
        IREE_STATUS_FAILED_PRECONDITION,
        "external replay file identity mismatch for '%.*s' captured as '%.*s': "
        "captured=(dev=%" PRIu64 ", inode=%" PRIu64 ", mtime_ns=%" PRIu64
        ") current=(dev=%" PRIu64 ", inode=%" PRIu64 ", mtime_ns=%" PRIu64
        "); identity validation is the default and does not read file "
        "contents; restore the original file identity, fix "
        "--replay_file_remap, or recapture with "
        "--device_replay_file_validation=digest when copied/staged files are "
        "intentional",
        (int)resolved_path.size, resolved_path.data, (int)captured_path.size,
        captured_path.data, payload->file_device, payload->file_inode,
        payload->file_mtime_ns, file_device, file_inode, file_mtime_ns);
  }
#else
  (void)handle;
  (void)captured_path;
  (void)resolved_path;
  if (payload->validation_type ==
      IREE_HAL_REPLAY_FILE_VALIDATION_TYPE_CONTENT_DIGEST) {
    return iree_make_status(
        IREE_STATUS_UNAVAILABLE,
        "external replay file content-digest validation requires POSIX file "
        "IO");
  }
#endif  // IREE_FILE_IO_ENABLE && (IREE_PLATFORM_ANDROID ||
        // IREE_PLATFORM_LINUX)
  return iree_ok_status();
}

static iree_status_t iree_hal_replay_executor_resolve_file_path(
    iree_hal_replay_executor_t* executor, iree_string_view_t captured_path,
    iree_string_view_t* out_resolved_path, char** out_resolved_path_storage) {
  IREE_ASSERT_ARGUMENT(out_resolved_path);
  IREE_ASSERT_ARGUMENT(out_resolved_path_storage);
  *out_resolved_path = captured_path;
  *out_resolved_path_storage = NULL;

  const iree_hal_replay_file_path_remap_t* selected_remap = NULL;
  for (iree_host_size_t i = 0; i < executor->options->file_path_remap_count;
       ++i) {
    const iree_hal_replay_file_path_remap_t* remap =
        &executor->options->file_path_remaps[i];
    if (iree_string_view_is_empty(remap->captured_prefix)) continue;
    if (!iree_string_view_starts_with(captured_path, remap->captured_prefix)) {
      continue;
    }
    if (!selected_remap ||
        remap->captured_prefix.size > selected_remap->captured_prefix.size) {
      selected_remap = remap;
    }
  }
  if (!selected_remap) return iree_ok_status();

  iree_string_view_t captured_suffix = iree_string_view_remove_prefix(
      captured_path, selected_remap->captured_prefix.size);
  iree_host_size_t resolved_length = 0;
  if (IREE_UNLIKELY(
          !iree_host_size_checked_add(selected_remap->replay_prefix.size,
                                      captured_suffix.size, &resolved_length) ||
          resolved_length == IREE_HOST_SIZE_MAX)) {
    return iree_make_status(IREE_STATUS_OUT_OF_RANGE,
                            "remapped replay file path is too long");
  }
  char* resolved_path_storage = NULL;
  IREE_RETURN_IF_ERROR(iree_allocator_malloc(executor->host_allocator,
                                             resolved_length + 1,
                                             (void**)&resolved_path_storage));
  memcpy(resolved_path_storage, selected_remap->replay_prefix.data,
         selected_remap->replay_prefix.size);
  memcpy(resolved_path_storage + selected_remap->replay_prefix.size,
         captured_suffix.data, captured_suffix.size);
  resolved_path_storage[resolved_length] = 0;
  *out_resolved_path =
      iree_make_string_view(resolved_path_storage, resolved_length);
  *out_resolved_path_storage = resolved_path_storage;
  return iree_ok_status();
}

typedef struct iree_hal_replay_executor_inline_file_release_t {
  // Host allocator used for the inline file copy and this release record.
  iree_allocator_t host_allocator;
} iree_hal_replay_executor_inline_file_release_t;

static void iree_hal_replay_executor_inline_file_release(
    void* user_data, iree_io_file_handle_primitive_t handle_primitive) {
  iree_hal_replay_executor_inline_file_release_t* release =
      (iree_hal_replay_executor_inline_file_release_t*)user_data;
  iree_allocator_t host_allocator = release->host_allocator;
  iree_allocator_free(host_allocator,
                      handle_primitive.value.host_allocation.data);
  iree_allocator_free(host_allocator, release);
}

static iree_io_file_access_t iree_hal_replay_executor_make_file_access(
    iree_hal_memory_access_t access) {
  iree_io_file_access_t file_access = 0;
  if (iree_any_bit_set(access, IREE_HAL_MEMORY_ACCESS_READ)) {
    file_access |= IREE_IO_FILE_ACCESS_READ;
  }
  if (iree_any_bit_set(access, IREE_HAL_MEMORY_ACCESS_WRITE)) {
    file_access |= IREE_IO_FILE_ACCESS_WRITE;
  }
  return file_access ? file_access : IREE_IO_FILE_ACCESS_READ;
}

static iree_status_t iree_hal_replay_executor_wrap_inline_file(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_file_object_payload_t* payload,
    iree_const_byte_span_t reference, iree_io_file_handle_t** out_handle) {
  IREE_ASSERT_ARGUMENT(out_handle);
  *out_handle = NULL;

  if (IREE_UNLIKELY(payload->file_length != reference.data_length)) {
    return iree_make_status(
        IREE_STATUS_DATA_LOSS,
        "inline replay file length mismatch: file_length=%" PRIu64
        " reference_length=%" PRIhsz,
        payload->file_length, reference.data_length);
  }

  iree_hal_replay_executor_inline_file_release_t* release = NULL;
  IREE_RETURN_IF_ERROR(iree_allocator_malloc(
      executor->host_allocator, sizeof(*release), (void**)&release));
  release->host_allocator = executor->host_allocator;

  uint8_t* file_bytes = NULL;
  iree_status_t status = iree_ok_status();
  if (reference.data_length != 0) {
    status = iree_allocator_malloc(executor->host_allocator,
                                   reference.data_length, (void**)&file_bytes);
    if (iree_status_is_ok(status)) {
      memcpy(file_bytes, reference.data, reference.data_length);
    }
  }
  iree_io_file_handle_release_callback_t release_callback = {
      .fn = iree_hal_replay_executor_inline_file_release,
      .user_data = release,
  };
  if (iree_status_is_ok(status)) {
    status = iree_io_file_handle_wrap_host_allocation(
        iree_hal_replay_executor_make_file_access(payload->access),
        iree_make_byte_span(file_bytes, reference.data_length),
        release_callback, executor->host_allocator, out_handle);
  }
  if (!iree_status_is_ok(status)) {
    iree_allocator_free(executor->host_allocator, file_bytes);
    iree_allocator_free(executor->host_allocator, release);
  }
  return status;
}

static iree_status_t iree_hal_replay_executor_import_file(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_file_record_t* record) {
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_require_payload(
      record, IREE_HAL_REPLAY_PAYLOAD_TYPE_FILE_OBJECT,
      sizeof(iree_hal_replay_file_object_payload_t)));
  iree_hal_replay_file_object_payload_t payload;
  memcpy(&payload, record->payload.data, sizeof(payload));
  if (IREE_UNLIKELY(payload.reference_length > IREE_HOST_SIZE_MAX ||
                    sizeof(payload) +
                            (iree_host_size_t)payload.reference_length !=
                        record->payload.data_length)) {
    return iree_make_status(IREE_STATUS_DATA_LOSS,
                            "replay file object payload length mismatch");
  }

  iree_hal_replay_object_entry_t* device_entry = NULL;
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_lookup(
      executor, record->header.object_id, IREE_HAL_REPLAY_OBJECT_TYPE_DEVICE,
      &device_entry));

  iree_const_byte_span_t reference =
      iree_make_const_byte_span(record->payload.data + sizeof(payload),
                                (iree_host_size_t)payload.reference_length);
  iree_io_file_handle_t* handle = NULL;
  char* resolved_path_storage = NULL;
  iree_string_view_t captured_external_path = iree_string_view_empty();
  iree_string_view_t resolved_external_path = iree_string_view_empty();
  bool import_file_handle = true;
  iree_status_t status = iree_ok_status();
  switch (payload.reference_type) {
    case IREE_HAL_REPLAY_FILE_REFERENCE_TYPE_EXTERNAL_PATH: {
      iree_string_view_t path = iree_make_string_view(
          (const char*)reference.data, reference.data_length);
      captured_external_path = path;
      status = iree_hal_replay_executor_resolve_file_path(
          executor, path, &path, &resolved_path_storage);
      resolved_external_path = path;
      iree_io_file_mode_t mode = 0;
      if (iree_any_bit_set(payload.access, IREE_HAL_MEMORY_ACCESS_READ)) {
        mode |= IREE_IO_FILE_MODE_READ;
      }
      if (iree_any_bit_set(payload.access, IREE_HAL_MEMORY_ACCESS_WRITE)) {
        mode |= IREE_IO_FILE_MODE_WRITE;
      }
      if (mode == 0) mode = IREE_IO_FILE_MODE_READ;
      if (iree_status_is_ok(status)) {
        status = iree_io_file_handle_open(mode, path, executor->host_allocator,
                                          &handle);
        if (!iree_status_is_ok(status)) {
          status = iree_status_annotate_f(
              status,
              "opening external replay file '%.*s' captured as '%.*s'; use "
              "--replay_file_remap=CAPTURED_PREFIX=REPLAY_PREFIX if the "
              "parameter root moved",
              (int)path.size, path.data, (int)captured_external_path.size,
              captured_external_path.data);
        }
      }
      break;
    }
    case IREE_HAL_REPLAY_FILE_REFERENCE_TYPE_INLINE_BYTES:
      status = iree_hal_replay_executor_wrap_inline_file(executor, &payload,
                                                         reference, &handle);
      break;
    case IREE_HAL_REPLAY_FILE_REFERENCE_TYPE_CAPTURED_RANGES:
      captured_external_path = iree_make_string_view(
          (const char*)reference.data, reference.data_length);
      import_file_handle = false;
      break;
    default:
      status = iree_make_status(IREE_STATUS_UNIMPLEMENTED,
                                "replay file reference type %" PRIu32
                                " is not executable",
                                payload.reference_type);
      break;
  }

  iree_hal_file_t* file = NULL;
  if (iree_status_is_ok(status) && import_file_handle) {
    status =
        iree_hal_file_import(device_entry->value.device, payload.queue_affinity,
                             payload.access, handle, payload.flags, &file);
  }
  if (iree_status_is_ok(status) &&
      payload.reference_type ==
          IREE_HAL_REPLAY_FILE_REFERENCE_TYPE_EXTERNAL_PATH) {
    status = iree_hal_replay_executor_validate_file_reference(
        &payload, captured_external_path, resolved_external_path, handle, file);
  }
  iree_io_file_handle_release(handle);
  iree_allocator_free(executor->host_allocator, resolved_path_storage);

  if (iree_status_is_ok(status)) {
    iree_hal_replay_object_entry_t entry = {.value.file = file};
    status = iree_hal_replay_executor_store(
        executor, record->header.related_object_id,
        IREE_HAL_REPLAY_OBJECT_TYPE_FILE, entry);
  } else {
    iree_hal_file_release(file);
  }
  return status;
}

static iree_status_t iree_hal_replay_executor_allocate_buffer(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_file_record_t* record) {
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_require_payload(
      record, IREE_HAL_REPLAY_PAYLOAD_TYPE_ALLOCATOR_ALLOCATE_BUFFER,
      sizeof(iree_hal_replay_allocator_allocate_buffer_payload_t)));
  iree_hal_replay_allocator_allocate_buffer_payload_t payload;
  memcpy(&payload, record->payload.data, sizeof(payload));
  iree_hal_replay_object_entry_t* allocator_entry = NULL;
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_lookup(
      executor, record->header.object_id, IREE_HAL_REPLAY_OBJECT_TYPE_ALLOCATOR,
      &allocator_entry));
  iree_hal_buffer_params_t params;
  IREE_RETURN_IF_ERROR(
      iree_hal_replay_executor_make_buffer_params(&payload, &params));
  iree_hal_buffer_t* buffer = NULL;
  IREE_RETURN_IF_ERROR(iree_hal_allocator_allocate_buffer(
      allocator_entry->value.allocator, params, payload.allocation_size,
      &buffer));
  iree_hal_replay_object_entry_t entry = {.value.buffer = buffer};
  return iree_hal_replay_executor_store(
      executor, record->header.related_object_id,
      IREE_HAL_REPLAY_OBJECT_TYPE_BUFFER, entry);
}

static iree_status_t iree_hal_replay_executor_import_buffer(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_file_record_t* record) {
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_require_payload(
      record, IREE_HAL_REPLAY_PAYLOAD_TYPE_ALLOCATOR_IMPORT_BUFFER,
      sizeof(iree_hal_replay_allocator_import_buffer_payload_t)));
  iree_hal_replay_allocator_import_buffer_payload_t payload;
  memcpy(&payload, record->payload.data, sizeof(payload));
  if (IREE_UNLIKELY(payload.external_type !=
                    IREE_HAL_EXTERNAL_BUFFER_TYPE_HOST_ALLOCATION)) {
    return iree_make_status(IREE_STATUS_UNIMPLEMENTED,
                            "replay import buffer external type %" PRIu32
                            " is not supported",
                            payload.external_type);
  }
  if (IREE_UNLIKELY(payload.data_length > IREE_HOST_SIZE_MAX ||
                    sizeof(payload) + (iree_host_size_t)payload.data_length !=
                        record->payload.data_length)) {
    return iree_make_status(IREE_STATUS_DATA_LOSS,
                            "replay import buffer payload length mismatch");
  }
  if (IREE_UNLIKELY(payload.data_length > payload.allocation.allocation_size)) {
    return iree_make_status(IREE_STATUS_DATA_LOSS,
                            "replay import buffer data overflows allocation");
  }
  iree_hal_replay_object_entry_t* allocator_entry = NULL;
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_lookup(
      executor, record->header.object_id, IREE_HAL_REPLAY_OBJECT_TYPE_ALLOCATOR,
      &allocator_entry));

  iree_hal_buffer_params_t params;
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_make_buffer_params(
      &payload.allocation, &params));
  params.type |= IREE_HAL_MEMORY_TYPE_HOST_VISIBLE;
  params.access |= IREE_HAL_MEMORY_ACCESS_WRITE;
  params.usage |= IREE_HAL_BUFFER_USAGE_MAPPING;

  iree_hal_buffer_t* buffer = NULL;
  iree_status_t status = iree_hal_allocator_allocate_buffer(
      allocator_entry->value.allocator, params,
      payload.allocation.allocation_size, &buffer);
  iree_const_byte_span_t data =
      iree_make_const_byte_span(record->payload.data + sizeof(payload),
                                (iree_host_size_t)payload.data_length);
  if (iree_status_is_ok(status)) {
    status = iree_hal_replay_executor_write_buffer_data(
        buffer, /*byte_offset=*/0, payload.allocation.allocation_size,
        IREE_HAL_MEMORY_ACCESS_WRITE, data);
  }
  if (iree_status_is_ok(status)) {
    iree_hal_replay_object_entry_t entry = {.value.buffer = buffer};
    status = iree_hal_replay_executor_store(
        executor, record->header.related_object_id,
        IREE_HAL_REPLAY_OBJECT_TYPE_BUFFER, entry);
  } else {
    iree_hal_buffer_release(buffer);
  }
  return status;
}

static iree_status_t iree_hal_replay_executor_replay_buffer_range_data(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_file_record_t* record) {
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_require_payload(
      record, IREE_HAL_REPLAY_PAYLOAD_TYPE_BUFFER_RANGE_DATA,
      sizeof(iree_hal_replay_buffer_range_data_payload_t)));
  iree_hal_replay_buffer_range_data_payload_t payload;
  memcpy(&payload, record->payload.data, sizeof(payload));
  if (IREE_UNLIKELY(payload.data_length >
                    record->payload.data_length - sizeof(payload))) {
    return iree_make_status(IREE_STATUS_DATA_LOSS,
                            "replay buffer range data overflows payload");
  }
  iree_hal_replay_object_entry_t* buffer_entry = NULL;
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_lookup(
      executor, record->header.object_id, IREE_HAL_REPLAY_OBJECT_TYPE_BUFFER,
      &buffer_entry));
  iree_const_byte_span_t data =
      iree_make_const_byte_span(record->payload.data + sizeof(payload),
                                (iree_host_size_t)payload.data_length);
  return iree_hal_replay_executor_write_buffer_data(
      buffer_entry->value.buffer, payload.byte_offset, payload.byte_length,
      payload.memory_access, data);
}

static iree_status_t iree_hal_replay_executor_queue_alloca(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_file_record_t* record) {
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_require_payload(
      record, IREE_HAL_REPLAY_PAYLOAD_TYPE_DEVICE_QUEUE_ALLOCA,
      sizeof(iree_hal_replay_device_queue_alloca_payload_t)));
  iree_hal_replay_device_queue_alloca_payload_t payload;
  memcpy(&payload, record->payload.data, sizeof(payload));

  iree_hal_replay_object_entry_t* device_entry = NULL;
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_lookup(
      executor, record->header.object_id, IREE_HAL_REPLAY_OBJECT_TYPE_DEVICE,
      &device_entry));
  iree_hal_replay_semaphore_list_storage_t wait_storage;
  iree_hal_replay_semaphore_list_storage_t signal_storage;
  iree_const_byte_span_t trailing_payload;
  iree_status_t status = iree_hal_replay_executor_make_queue_semaphore_lists(
      executor, record, sizeof(payload), payload.wait_semaphore_count,
      payload.signal_semaphore_count, /*trailing_payload_length=*/0,
      &wait_storage, &signal_storage, &trailing_payload);
  iree_hal_buffer_t* buffer = NULL;
  if (iree_status_is_ok(status)) {
    iree_hal_buffer_params_t params;
    status = iree_hal_replay_executor_make_buffer_params(&payload.allocation,
                                                         &params);
    if (iree_status_is_ok(status)) {
      status = iree_hal_device_queue_alloca(
          device_entry->value.device, payload.queue_affinity, wait_storage.list,
          signal_storage.list, /*pool=*/NULL, params,
          payload.allocation.allocation_size, payload.flags, &buffer);
    }
  }
  if (iree_status_is_ok(status)) {
    status = iree_hal_replay_executor_flush_and_wait(device_entry->value.device,
                                                     payload.queue_affinity,
                                                     signal_storage.list);
  }
  if (iree_status_is_ok(status)) {
    iree_hal_replay_object_entry_t entry = {.value.buffer = buffer};
    buffer = NULL;
    status = iree_hal_replay_executor_store(
        executor, record->header.related_object_id,
        IREE_HAL_REPLAY_OBJECT_TYPE_BUFFER, entry);
  }
  iree_hal_buffer_release(buffer);
  iree_hal_replay_semaphore_list_storage_deinitialize(&signal_storage,
                                                      executor->host_allocator);
  iree_hal_replay_semaphore_list_storage_deinitialize(&wait_storage,
                                                      executor->host_allocator);
  return status;
}

static iree_status_t iree_hal_replay_executor_queue_dealloca(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_file_record_t* record) {
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_require_payload(
      record, IREE_HAL_REPLAY_PAYLOAD_TYPE_DEVICE_QUEUE_DEALLOCA,
      sizeof(iree_hal_replay_device_queue_dealloca_payload_t)));
  iree_hal_replay_device_queue_dealloca_payload_t payload;
  memcpy(&payload, record->payload.data, sizeof(payload));

  iree_hal_replay_object_entry_t* device_entry = NULL;
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_lookup(
      executor, record->header.object_id, IREE_HAL_REPLAY_OBJECT_TYPE_DEVICE,
      &device_entry));
  iree_hal_replay_semaphore_list_storage_t wait_storage;
  iree_hal_replay_semaphore_list_storage_t signal_storage;
  iree_const_byte_span_t trailing_payload;
  iree_status_t status = iree_hal_replay_executor_make_queue_semaphore_lists(
      executor, record, sizeof(payload), payload.wait_semaphore_count,
      payload.signal_semaphore_count, /*trailing_payload_length=*/0,
      &wait_storage, &signal_storage, &trailing_payload);
  iree_hal_buffer_ref_t buffer_ref;
  if (iree_status_is_ok(status)) {
    status = iree_hal_replay_executor_make_buffer_ref(
        executor, &payload.buffer_ref, &buffer_ref);
  }
  if (iree_status_is_ok(status) && !buffer_ref.buffer) {
    status = iree_make_status(
        IREE_STATUS_DATA_LOSS,
        "replay queue dealloca requires a direct buffer reference");
  }
  if (iree_status_is_ok(status)) {
    status = iree_hal_device_queue_dealloca(
        device_entry->value.device, payload.queue_affinity, wait_storage.list,
        signal_storage.list, buffer_ref.buffer, payload.flags);
  }
  if (iree_status_is_ok(status)) {
    status = iree_hal_replay_executor_flush_and_wait(device_entry->value.device,
                                                     payload.queue_affinity,
                                                     signal_storage.list);
  }
  iree_hal_replay_semaphore_list_storage_deinitialize(&signal_storage,
                                                      executor->host_allocator);
  iree_hal_replay_semaphore_list_storage_deinitialize(&wait_storage,
                                                      executor->host_allocator);
  return status;
}

static iree_status_t iree_hal_replay_executor_queue_fill(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_file_record_t* record) {
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_require_payload(
      record, IREE_HAL_REPLAY_PAYLOAD_TYPE_DEVICE_QUEUE_FILL,
      sizeof(iree_hal_replay_device_queue_fill_payload_t)));
  iree_hal_replay_device_queue_fill_payload_t payload;
  memcpy(&payload, record->payload.data, sizeof(payload));

  iree_hal_replay_object_entry_t* device_entry = NULL;
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_lookup(
      executor, record->header.object_id, IREE_HAL_REPLAY_OBJECT_TYPE_DEVICE,
      &device_entry));
  iree_hal_replay_semaphore_list_storage_t wait_storage;
  iree_hal_replay_semaphore_list_storage_t signal_storage;
  iree_const_byte_span_t pattern;
  iree_status_t status = iree_hal_replay_executor_make_queue_semaphore_lists(
      executor, record, sizeof(payload), payload.wait_semaphore_count,
      payload.signal_semaphore_count, payload.pattern_length, &wait_storage,
      &signal_storage, &pattern);
  iree_hal_buffer_ref_t target_ref;
  if (iree_status_is_ok(status)) {
    status = iree_hal_replay_executor_make_buffer_ref(
        executor, &payload.target_ref, &target_ref);
  }
  if (iree_status_is_ok(status) && !target_ref.buffer) {
    status = iree_make_status(
        IREE_STATUS_DATA_LOSS,
        "replay queue fill requires a direct target buffer reference");
  }
  if (iree_status_is_ok(status)) {
    status = iree_hal_device_queue_fill(
        device_entry->value.device, payload.queue_affinity, wait_storage.list,
        signal_storage.list, target_ref.buffer, target_ref.offset,
        target_ref.length, pattern.data, pattern.data_length, payload.flags);
  }
  if (iree_status_is_ok(status)) {
    status = iree_hal_replay_executor_flush_and_wait(device_entry->value.device,
                                                     payload.queue_affinity,
                                                     signal_storage.list);
  }
  iree_hal_replay_semaphore_list_storage_deinitialize(&signal_storage,
                                                      executor->host_allocator);
  iree_hal_replay_semaphore_list_storage_deinitialize(&wait_storage,
                                                      executor->host_allocator);
  return status;
}

static iree_status_t iree_hal_replay_executor_queue_update(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_file_record_t* record) {
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_require_payload(
      record, IREE_HAL_REPLAY_PAYLOAD_TYPE_DEVICE_QUEUE_UPDATE,
      sizeof(iree_hal_replay_device_queue_update_payload_t)));
  iree_hal_replay_device_queue_update_payload_t payload;
  memcpy(&payload, record->payload.data, sizeof(payload));

  iree_hal_replay_object_entry_t* device_entry = NULL;
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_lookup(
      executor, record->header.object_id, IREE_HAL_REPLAY_OBJECT_TYPE_DEVICE,
      &device_entry));
  iree_hal_replay_semaphore_list_storage_t wait_storage;
  iree_hal_replay_semaphore_list_storage_t signal_storage;
  iree_const_byte_span_t data;
  iree_status_t status = iree_hal_replay_executor_make_queue_semaphore_lists(
      executor, record, sizeof(payload), payload.wait_semaphore_count,
      payload.signal_semaphore_count, payload.data_length, &wait_storage,
      &signal_storage, &data);
  iree_hal_buffer_ref_t target_ref;
  if (iree_status_is_ok(status)) {
    status = iree_hal_replay_executor_make_buffer_ref(
        executor, &payload.target_ref, &target_ref);
  }
  if (iree_status_is_ok(status) && !target_ref.buffer) {
    status = iree_make_status(
        IREE_STATUS_DATA_LOSS,
        "replay queue update requires a direct target buffer reference");
  }
  if (iree_status_is_ok(status) && data.data_length != target_ref.length) {
    status = iree_make_status(
        IREE_STATUS_DATA_LOSS,
        "replay queue update data length does not match target length");
  }
  if (iree_status_is_ok(status)) {
    status = iree_hal_device_queue_update(
        device_entry->value.device, payload.queue_affinity, wait_storage.list,
        signal_storage.list, data.data, /*source_offset=*/0, target_ref.buffer,
        target_ref.offset, target_ref.length, payload.flags);
  }
  if (iree_status_is_ok(status)) {
    status = iree_hal_replay_executor_flush_and_wait(device_entry->value.device,
                                                     payload.queue_affinity,
                                                     signal_storage.list);
  }
  iree_hal_replay_semaphore_list_storage_deinitialize(&signal_storage,
                                                      executor->host_allocator);
  iree_hal_replay_semaphore_list_storage_deinitialize(&wait_storage,
                                                      executor->host_allocator);
  return status;
}

static iree_status_t iree_hal_replay_executor_queue_copy(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_file_record_t* record) {
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_require_payload(
      record, IREE_HAL_REPLAY_PAYLOAD_TYPE_DEVICE_QUEUE_COPY,
      sizeof(iree_hal_replay_device_queue_copy_payload_t)));
  iree_hal_replay_device_queue_copy_payload_t payload;
  memcpy(&payload, record->payload.data, sizeof(payload));

  iree_hal_replay_object_entry_t* device_entry = NULL;
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_lookup(
      executor, record->header.object_id, IREE_HAL_REPLAY_OBJECT_TYPE_DEVICE,
      &device_entry));
  iree_hal_replay_semaphore_list_storage_t wait_storage;
  iree_hal_replay_semaphore_list_storage_t signal_storage;
  iree_const_byte_span_t trailing_payload;
  iree_status_t status = iree_hal_replay_executor_make_queue_semaphore_lists(
      executor, record, sizeof(payload), payload.wait_semaphore_count,
      payload.signal_semaphore_count, /*trailing_payload_length=*/0,
      &wait_storage, &signal_storage, &trailing_payload);
  iree_hal_buffer_ref_t source_ref;
  if (iree_status_is_ok(status)) {
    status = iree_hal_replay_executor_make_buffer_ref(
        executor, &payload.source_ref, &source_ref);
  }
  iree_hal_buffer_ref_t target_ref;
  if (iree_status_is_ok(status)) {
    status = iree_hal_replay_executor_make_buffer_ref(
        executor, &payload.target_ref, &target_ref);
  }
  if (iree_status_is_ok(status) && (!source_ref.buffer || !target_ref.buffer)) {
    status = iree_make_status(IREE_STATUS_DATA_LOSS,
                              "replay queue copy requires direct source and "
                              "target buffer references");
  }
  if (iree_status_is_ok(status) && source_ref.length != target_ref.length) {
    status = iree_make_status(
        IREE_STATUS_DATA_LOSS,
        "replay queue copy source and target lengths do not match");
  }
  if (iree_status_is_ok(status)) {
    status = iree_hal_device_queue_copy(
        device_entry->value.device, payload.queue_affinity, wait_storage.list,
        signal_storage.list, source_ref.buffer, source_ref.offset,
        target_ref.buffer, target_ref.offset, target_ref.length, payload.flags);
  }
  if (iree_status_is_ok(status)) {
    status = iree_hal_replay_executor_flush_and_wait(device_entry->value.device,
                                                     payload.queue_affinity,
                                                     signal_storage.list);
  }
  iree_hal_replay_semaphore_list_storage_deinitialize(&signal_storage,
                                                      executor->host_allocator);
  iree_hal_replay_semaphore_list_storage_deinitialize(&wait_storage,
                                                      executor->host_allocator);
  return status;
}

static iree_status_t iree_hal_replay_executor_queue_read(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_file_record_t* record) {
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_require_payload(
      record, IREE_HAL_REPLAY_PAYLOAD_TYPE_DEVICE_QUEUE_READ,
      sizeof(iree_hal_replay_device_queue_read_payload_t)));
  iree_hal_replay_device_queue_read_payload_t payload;
  memcpy(&payload, record->payload.data, sizeof(payload));

  iree_hal_replay_object_entry_t* device_entry = NULL;
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_lookup(
      executor, record->header.object_id, IREE_HAL_REPLAY_OBJECT_TYPE_DEVICE,
      &device_entry));
  iree_hal_replay_semaphore_list_storage_t wait_storage;
  iree_hal_replay_semaphore_list_storage_t signal_storage;
  iree_const_byte_span_t trailing_payload;
  iree_status_t status = iree_hal_replay_executor_make_queue_semaphore_lists(
      executor, record, sizeof(payload), payload.wait_semaphore_count,
      payload.signal_semaphore_count, payload.captured_data_length,
      &wait_storage, &signal_storage, &trailing_payload);
  iree_hal_replay_object_entry_t* file_entry = NULL;
  if (iree_status_is_ok(status)) {
    status = iree_hal_replay_executor_lookup(executor, payload.source_file_id,
                                             IREE_HAL_REPLAY_OBJECT_TYPE_FILE,
                                             &file_entry);
  }
  iree_hal_buffer_ref_t target_ref;
  if (iree_status_is_ok(status)) {
    status = iree_hal_replay_executor_make_buffer_ref(
        executor, &payload.target_ref, &target_ref);
  }
  if (iree_status_is_ok(status) && !target_ref.buffer) {
    status = iree_make_status(
        IREE_STATUS_DATA_LOSS,
        "replay queue read requires a direct target buffer reference");
  }
  if (iree_status_is_ok(status) && payload.captured_data_length != 0 &&
      payload.captured_data_length != target_ref.length) {
    status = iree_make_status(
        IREE_STATUS_DATA_LOSS,
        "replay queue read captured data length does not match target length");
  }
  if (iree_status_is_ok(status) && payload.captured_data_length != 0) {
    status = iree_hal_device_queue_update(
        device_entry->value.device, payload.queue_affinity, wait_storage.list,
        signal_storage.list, trailing_payload.data, /*source_offset=*/0,
        target_ref.buffer, target_ref.offset, target_ref.length,
        IREE_HAL_UPDATE_FLAG_NONE);
  } else if (iree_status_is_ok(status)) {
    if (!file_entry->value.file) {
      status = iree_make_status(
          IREE_STATUS_FAILED_PRECONDITION,
          "replay queue read requires an imported file or captured data");
    }
  }
  if (iree_status_is_ok(status) && payload.captured_data_length == 0) {
    status = iree_hal_device_queue_read(
        device_entry->value.device, payload.queue_affinity, wait_storage.list,
        signal_storage.list, file_entry->value.file, payload.source_offset,
        target_ref.buffer, target_ref.offset, target_ref.length, payload.flags);
  }
  if (iree_status_is_ok(status)) {
    status = iree_hal_replay_executor_flush_and_wait(device_entry->value.device,
                                                     payload.queue_affinity,
                                                     signal_storage.list);
  }
  iree_hal_replay_semaphore_list_storage_deinitialize(&signal_storage,
                                                      executor->host_allocator);
  iree_hal_replay_semaphore_list_storage_deinitialize(&wait_storage,
                                                      executor->host_allocator);
  return status;
}

static iree_status_t iree_hal_replay_executor_queue_write(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_file_record_t* record) {
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_require_payload(
      record, IREE_HAL_REPLAY_PAYLOAD_TYPE_DEVICE_QUEUE_WRITE,
      sizeof(iree_hal_replay_device_queue_write_payload_t)));
  iree_hal_replay_device_queue_write_payload_t payload;
  memcpy(&payload, record->payload.data, sizeof(payload));

  iree_hal_replay_object_entry_t* device_entry = NULL;
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_lookup(
      executor, record->header.object_id, IREE_HAL_REPLAY_OBJECT_TYPE_DEVICE,
      &device_entry));
  iree_hal_replay_semaphore_list_storage_t wait_storage;
  iree_hal_replay_semaphore_list_storage_t signal_storage;
  iree_const_byte_span_t trailing_payload;
  iree_status_t status = iree_hal_replay_executor_make_queue_semaphore_lists(
      executor, record, sizeof(payload), payload.wait_semaphore_count,
      payload.signal_semaphore_count, /*trailing_payload_length=*/0,
      &wait_storage, &signal_storage, &trailing_payload);
  iree_hal_buffer_ref_t source_ref;
  if (iree_status_is_ok(status)) {
    status = iree_hal_replay_executor_make_buffer_ref(
        executor, &payload.source_ref, &source_ref);
  }
  iree_hal_replay_object_entry_t* file_entry = NULL;
  if (iree_status_is_ok(status)) {
    status = iree_hal_replay_executor_lookup(executor, payload.target_file_id,
                                             IREE_HAL_REPLAY_OBJECT_TYPE_FILE,
                                             &file_entry);
  }
  if (iree_status_is_ok(status) && !source_ref.buffer) {
    status = iree_make_status(
        IREE_STATUS_DATA_LOSS,
        "replay queue write requires a direct source buffer reference");
  }
  if (iree_status_is_ok(status) && !file_entry->value.file) {
    status =
        iree_make_status(IREE_STATUS_FAILED_PRECONDITION,
                         "replay queue write requires an imported target file");
  }
  if (iree_status_is_ok(status)) {
    status = iree_hal_device_queue_write(
        device_entry->value.device, payload.queue_affinity, wait_storage.list,
        signal_storage.list, source_ref.buffer, source_ref.offset,
        file_entry->value.file, payload.target_offset, source_ref.length,
        payload.flags);
  }
  if (iree_status_is_ok(status)) {
    status = iree_hal_replay_executor_flush_and_wait(device_entry->value.device,
                                                     payload.queue_affinity,
                                                     signal_storage.list);
  }
  iree_hal_replay_semaphore_list_storage_deinitialize(&signal_storage,
                                                      executor->host_allocator);
  iree_hal_replay_semaphore_list_storage_deinitialize(&wait_storage,
                                                      executor->host_allocator);
  return status;
}

static iree_status_t iree_hal_replay_executor_command_buffer_execution_barrier(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_file_record_t* record) {
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_require_payload(
      record, IREE_HAL_REPLAY_PAYLOAD_TYPE_COMMAND_BUFFER_EXECUTION_BARRIER,
      sizeof(iree_hal_replay_command_buffer_execution_barrier_payload_t)));
  iree_hal_replay_command_buffer_execution_barrier_payload_t payload;
  memcpy(&payload, record->payload.data, sizeof(payload));
  iree_host_size_t memory_payloads_size = 0;
  iree_host_size_t buffer_payloads_size = 0;
  iree_host_size_t total_payload_size = 0;
  if (IREE_UNLIKELY(
          payload.memory_barrier_count > IREE_HOST_SIZE_MAX ||
          payload.buffer_barrier_count > IREE_HOST_SIZE_MAX ||
          !iree_host_size_checked_mul(
              (iree_host_size_t)payload.memory_barrier_count,
              sizeof(iree_hal_replay_memory_barrier_payload_t),
              &memory_payloads_size) ||
          !iree_host_size_checked_mul(
              (iree_host_size_t)payload.buffer_barrier_count,
              sizeof(iree_hal_replay_buffer_barrier_payload_t),
              &buffer_payloads_size) ||
          !iree_host_size_checked_add(sizeof(payload), memory_payloads_size,
                                      &total_payload_size) ||
          !iree_host_size_checked_add(total_payload_size, buffer_payloads_size,
                                      &total_payload_size) ||
          total_payload_size != record->payload.data_length)) {
    return iree_make_status(IREE_STATUS_DATA_LOSS,
                            "replay execution barrier payload length mismatch");
  }

  iree_hal_memory_barrier_t inline_memory_barriers
      [IREE_HAL_REPLAY_INLINE_MEMORY_BARRIER_LIST_CAPACITY];
  iree_hal_memory_barrier_t* memory_barriers = NULL;
  bool memory_barriers_allocated = false;
  if (payload.memory_barrier_count <=
      IREE_HAL_REPLAY_INLINE_MEMORY_BARRIER_LIST_CAPACITY) {
    memory_barriers = inline_memory_barriers;
  } else {
    iree_host_size_t memory_barriers_size = 0;
    if (IREE_UNLIKELY(!iree_host_size_checked_mul(
            (iree_host_size_t)payload.memory_barrier_count,
            sizeof(*memory_barriers), &memory_barriers_size))) {
      return iree_make_status(
          IREE_STATUS_OUT_OF_RANGE,
          "replay execution barrier memory barrier count overflow");
    }
    IREE_RETURN_IF_ERROR(iree_allocator_malloc(executor->host_allocator,
                                               memory_barriers_size,
                                               (void**)&memory_barriers));
    memory_barriers_allocated = true;
  }
  const iree_hal_replay_memory_barrier_payload_t* memory_payloads =
      (const iree_hal_replay_memory_barrier_payload_t*)(record->payload.data +
                                                        sizeof(payload));
  for (iree_host_size_t i = 0; i < payload.memory_barrier_count; ++i) {
    memory_barriers[i].source_scope = memory_payloads[i].source_scope;
    memory_barriers[i].target_scope = memory_payloads[i].target_scope;
  }

  iree_hal_buffer_barrier_t inline_buffer_barriers
      [IREE_HAL_REPLAY_INLINE_BUFFER_BARRIER_LIST_CAPACITY];
  iree_hal_buffer_barrier_t* buffer_barriers = NULL;
  iree_status_t status = iree_ok_status();
  bool buffer_barriers_allocated = false;
  if (payload.buffer_barrier_count <=
      IREE_HAL_REPLAY_INLINE_BUFFER_BARRIER_LIST_CAPACITY) {
    buffer_barriers = inline_buffer_barriers;
  } else {
    iree_host_size_t buffer_barriers_size = 0;
    if (IREE_UNLIKELY(!iree_host_size_checked_mul(
            (iree_host_size_t)payload.buffer_barrier_count,
            sizeof(*buffer_barriers), &buffer_barriers_size))) {
      status = iree_make_status(
          IREE_STATUS_OUT_OF_RANGE,
          "replay execution barrier buffer barrier count overflow");
    }
    if (iree_status_is_ok(status)) {
      status =
          iree_allocator_malloc(executor->host_allocator, buffer_barriers_size,
                                (void**)&buffer_barriers);
    }
    buffer_barriers_allocated = iree_status_is_ok(status);
  }
  const iree_hal_replay_buffer_barrier_payload_t* buffer_payloads =
      (const iree_hal_replay_buffer_barrier_payload_t*)(record->payload.data +
                                                        sizeof(payload) +
                                                        memory_payloads_size);
  for (iree_host_size_t i = 0;
       i < payload.buffer_barrier_count && iree_status_is_ok(status); ++i) {
    buffer_barriers[i].source_scope = buffer_payloads[i].source_scope;
    buffer_barriers[i].target_scope = buffer_payloads[i].target_scope;
    status = iree_hal_replay_executor_make_buffer_ref(
        executor, &buffer_payloads[i].buffer_ref,
        &buffer_barriers[i].buffer_ref);
  }

  iree_hal_replay_object_entry_t* command_buffer_entry = NULL;
  if (iree_status_is_ok(status)) {
    status = iree_hal_replay_executor_lookup(
        executor, record->header.object_id,
        IREE_HAL_REPLAY_OBJECT_TYPE_COMMAND_BUFFER, &command_buffer_entry);
  }
  if (iree_status_is_ok(status)) {
    status = iree_hal_command_buffer_execution_barrier(
        command_buffer_entry->value.command_buffer, payload.source_stage_mask,
        payload.target_stage_mask, payload.flags,
        (iree_host_size_t)payload.memory_barrier_count, memory_barriers,
        (iree_host_size_t)payload.buffer_barrier_count, buffer_barriers);
  }
  if (buffer_barriers_allocated) {
    iree_allocator_free(executor->host_allocator, buffer_barriers);
  }
  if (memory_barriers_allocated) {
    iree_allocator_free(executor->host_allocator, memory_barriers);
  }
  return status;
}

static iree_status_t iree_hal_replay_executor_command_buffer_event(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_file_record_t* record) {
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_require_payload(
      record, IREE_HAL_REPLAY_PAYLOAD_TYPE_COMMAND_BUFFER_EVENT,
      sizeof(iree_hal_replay_command_buffer_event_payload_t)));
  iree_hal_replay_command_buffer_event_payload_t payload;
  memcpy(&payload, record->payload.data, sizeof(payload));

  iree_hal_replay_object_entry_t* command_buffer_entry = NULL;
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_lookup(
      executor, record->header.object_id,
      IREE_HAL_REPLAY_OBJECT_TYPE_COMMAND_BUFFER, &command_buffer_entry));
  iree_hal_replay_object_entry_t* event_entry = NULL;
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_lookup(
      executor, payload.event_id, IREE_HAL_REPLAY_OBJECT_TYPE_EVENT,
      &event_entry));

  switch (record->header.operation_code) {
    case IREE_HAL_REPLAY_OPERATION_CODE_COMMAND_BUFFER_SIGNAL_EVENT:
      return iree_hal_command_buffer_signal_event(
          command_buffer_entry->value.command_buffer, event_entry->value.event,
          payload.source_stage_mask);
    case IREE_HAL_REPLAY_OPERATION_CODE_COMMAND_BUFFER_RESET_EVENT:
      return iree_hal_command_buffer_reset_event(
          command_buffer_entry->value.command_buffer, event_entry->value.event,
          payload.source_stage_mask);
    default:
      return iree_make_status(
          IREE_STATUS_UNIMPLEMENTED, "replay operation %s is not an event op",
          iree_hal_replay_operation_code_string(record->header.operation_code));
  }
}

static iree_status_t iree_hal_replay_executor_command_buffer_wait_events(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_file_record_t* record) {
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_require_payload(
      record, IREE_HAL_REPLAY_PAYLOAD_TYPE_COMMAND_BUFFER_WAIT_EVENTS,
      sizeof(iree_hal_replay_command_buffer_wait_events_payload_t)));
  iree_hal_replay_command_buffer_wait_events_payload_t payload;
  memcpy(&payload, record->payload.data, sizeof(payload));

  iree_host_size_t event_payloads_size = 0;
  iree_host_size_t memory_payloads_size = 0;
  iree_host_size_t buffer_payloads_size = 0;
  iree_host_size_t total_payload_size = 0;
  if (IREE_UNLIKELY(
          payload.event_count > IREE_HOST_SIZE_MAX ||
          payload.memory_barrier_count > IREE_HOST_SIZE_MAX ||
          payload.buffer_barrier_count > IREE_HOST_SIZE_MAX ||
          !iree_host_size_checked_mul((iree_host_size_t)payload.event_count,
                                      sizeof(iree_hal_replay_object_id_t),
                                      &event_payloads_size) ||
          !iree_host_size_checked_mul(
              (iree_host_size_t)payload.memory_barrier_count,
              sizeof(iree_hal_replay_memory_barrier_payload_t),
              &memory_payloads_size) ||
          !iree_host_size_checked_mul(
              (iree_host_size_t)payload.buffer_barrier_count,
              sizeof(iree_hal_replay_buffer_barrier_payload_t),
              &buffer_payloads_size) ||
          !iree_host_size_checked_add(sizeof(payload), event_payloads_size,
                                      &total_payload_size) ||
          !iree_host_size_checked_add(total_payload_size, memory_payloads_size,
                                      &total_payload_size) ||
          !iree_host_size_checked_add(total_payload_size, buffer_payloads_size,
                                      &total_payload_size) ||
          total_payload_size != record->payload.data_length)) {
    return iree_make_status(IREE_STATUS_DATA_LOSS,
                            "replay wait events payload length mismatch");
  }

  const iree_hal_replay_object_id_t* event_payloads =
      (const iree_hal_replay_object_id_t*)(record->payload.data +
                                           sizeof(payload));
  const iree_hal_replay_memory_barrier_payload_t* memory_payloads =
      (const iree_hal_replay_memory_barrier_payload_t*)(record->payload.data +
                                                        sizeof(payload) +
                                                        event_payloads_size);
  const iree_hal_replay_buffer_barrier_payload_t* buffer_payloads =
      (const iree_hal_replay_buffer_barrier_payload_t*)(record->payload.data +
                                                        sizeof(payload) +
                                                        event_payloads_size +
                                                        memory_payloads_size);

  const iree_hal_event_t** events = NULL;
  iree_status_t status = iree_ok_status();
  iree_host_size_t events_size = 0;
  if (payload.event_count) {
    if (IREE_UNLIKELY(
            !iree_host_size_checked_mul((iree_host_size_t)payload.event_count,
                                        sizeof(*events), &events_size))) {
      status = iree_make_status(IREE_STATUS_OUT_OF_RANGE,
                                "replay wait event count overflow");
    }
    if (iree_status_is_ok(status)) {
      status = iree_allocator_malloc(executor->host_allocator, events_size,
                                     (void**)&events);
    }
  }
  for (iree_host_size_t i = 0;
       i < payload.event_count && iree_status_is_ok(status); ++i) {
    iree_hal_replay_object_entry_t* event_entry = NULL;
    status = iree_hal_replay_executor_lookup(executor, event_payloads[i],
                                             IREE_HAL_REPLAY_OBJECT_TYPE_EVENT,
                                             &event_entry);
    if (iree_status_is_ok(status)) {
      events[i] = event_entry->value.event;
    }
  }

  iree_hal_memory_barrier_t* memory_barriers = NULL;
  iree_host_size_t memory_barriers_size = 0;
  if (iree_status_is_ok(status) && payload.memory_barrier_count) {
    if (IREE_UNLIKELY(!iree_host_size_checked_mul(
            (iree_host_size_t)payload.memory_barrier_count,
            sizeof(*memory_barriers), &memory_barriers_size))) {
      status = iree_make_status(IREE_STATUS_OUT_OF_RANGE,
                                "replay wait memory barrier count overflow");
    }
    if (iree_status_is_ok(status)) {
      status =
          iree_allocator_malloc(executor->host_allocator, memory_barriers_size,
                                (void**)&memory_barriers);
    }
  }
  for (iree_host_size_t i = 0;
       i < payload.memory_barrier_count && iree_status_is_ok(status); ++i) {
    memory_barriers[i].source_scope = memory_payloads[i].source_scope;
    memory_barriers[i].target_scope = memory_payloads[i].target_scope;
  }

  iree_hal_buffer_barrier_t* buffer_barriers = NULL;
  iree_host_size_t buffer_barriers_size = 0;
  if (iree_status_is_ok(status) && payload.buffer_barrier_count) {
    if (IREE_UNLIKELY(!iree_host_size_checked_mul(
            (iree_host_size_t)payload.buffer_barrier_count,
            sizeof(*buffer_barriers), &buffer_barriers_size))) {
      status = iree_make_status(IREE_STATUS_OUT_OF_RANGE,
                                "replay wait buffer barrier count overflow");
    }
    if (iree_status_is_ok(status)) {
      status =
          iree_allocator_malloc(executor->host_allocator, buffer_barriers_size,
                                (void**)&buffer_barriers);
    }
  }
  for (iree_host_size_t i = 0;
       i < payload.buffer_barrier_count && iree_status_is_ok(status); ++i) {
    buffer_barriers[i].source_scope = buffer_payloads[i].source_scope;
    buffer_barriers[i].target_scope = buffer_payloads[i].target_scope;
    status = iree_hal_replay_executor_make_buffer_ref(
        executor, &buffer_payloads[i].buffer_ref,
        &buffer_barriers[i].buffer_ref);
  }

  iree_hal_replay_object_entry_t* command_buffer_entry = NULL;
  if (iree_status_is_ok(status)) {
    status = iree_hal_replay_executor_lookup(
        executor, record->header.object_id,
        IREE_HAL_REPLAY_OBJECT_TYPE_COMMAND_BUFFER, &command_buffer_entry);
  }
  if (iree_status_is_ok(status)) {
    status = iree_hal_command_buffer_wait_events(
        command_buffer_entry->value.command_buffer,
        (iree_host_size_t)payload.event_count, events,
        payload.source_stage_mask, payload.target_stage_mask,
        (iree_host_size_t)payload.memory_barrier_count, memory_barriers,
        (iree_host_size_t)payload.buffer_barrier_count, buffer_barriers);
  }

  iree_allocator_free(executor->host_allocator, buffer_barriers);
  iree_allocator_free(executor->host_allocator, memory_barriers);
  iree_allocator_free(executor->host_allocator, events);
  return status;
}

static iree_status_t iree_hal_replay_executor_dispatch_layout(
    const iree_hal_replay_file_record_t* record,
    const iree_hal_replay_dispatch_payload_t* payload,
    iree_host_size_t* out_wait_payloads_offset,
    iree_host_size_t* out_wait_payloads_size,
    iree_host_size_t* out_signal_payloads_offset,
    iree_host_size_t* out_signal_payloads_size,
    iree_host_size_t* out_constants_offset,
    iree_host_size_t* out_binding_payloads_offset,
    iree_host_size_t* out_binding_payloads_size) {
  iree_host_size_t wait_size = 0;
  iree_host_size_t signal_size = 0;
  iree_host_size_t constants_size = 0;
  iree_host_size_t binding_size = 0;
  if (IREE_UNLIKELY(payload->wait_semaphore_count > IREE_HOST_SIZE_MAX ||
                    payload->signal_semaphore_count > IREE_HOST_SIZE_MAX ||
                    payload->constants_length > IREE_HOST_SIZE_MAX ||
                    payload->binding_count > IREE_HOST_SIZE_MAX ||
                    !iree_host_size_checked_mul(
                        (iree_host_size_t)payload->wait_semaphore_count,
                        sizeof(iree_hal_replay_semaphore_timepoint_payload_t),
                        &wait_size) ||
                    !iree_host_size_checked_mul(
                        (iree_host_size_t)payload->signal_semaphore_count,
                        sizeof(iree_hal_replay_semaphore_timepoint_payload_t),
                        &signal_size) ||
                    !iree_host_size_checked_mul(
                        (iree_host_size_t)payload->binding_count,
                        sizeof(iree_hal_replay_buffer_ref_payload_t),
                        &binding_size))) {
    return iree_make_status(IREE_STATUS_OUT_OF_RANGE,
                            "replay dispatch payload count overflow");
  }
  constants_size = (iree_host_size_t)payload->constants_length;
  iree_host_size_t offset = sizeof(*payload);
  *out_wait_payloads_offset = offset;
  *out_wait_payloads_size = wait_size;
  if (!iree_host_size_checked_add(offset, wait_size, &offset)) {
    return iree_make_status(IREE_STATUS_OUT_OF_RANGE,
                            "replay dispatch payload length overflow");
  }
  *out_signal_payloads_offset = offset;
  *out_signal_payloads_size = signal_size;
  if (!iree_host_size_checked_add(offset, signal_size, &offset)) {
    return iree_make_status(IREE_STATUS_OUT_OF_RANGE,
                            "replay dispatch payload length overflow");
  }
  *out_constants_offset = offset;
  if (!iree_host_size_checked_add(offset, constants_size, &offset)) {
    return iree_make_status(IREE_STATUS_OUT_OF_RANGE,
                            "replay dispatch payload length overflow");
  }
  *out_binding_payloads_offset = offset;
  *out_binding_payloads_size = binding_size;
  if (!iree_host_size_checked_add(offset, binding_size, &offset) ||
      offset != record->payload.data_length) {
    return iree_make_status(IREE_STATUS_DATA_LOSS,
                            "replay dispatch payload length mismatch");
  }
  return iree_ok_status();
}

static iree_status_t iree_hal_replay_plan_prepare_scope(
    const iree_hal_replay_file_record_t* record,
    iree_hal_replay_scope_event_type_t event_type,
    iree_hal_replay_plan_scope_t* out_scope) {
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_require_payload(
      record, IREE_HAL_REPLAY_PAYLOAD_TYPE_REPLAY_SCOPE,
      sizeof(iree_hal_replay_scope_payload_t)));
  iree_hal_replay_scope_payload_t payload;
  memcpy(&payload, record->payload.data, sizeof(payload));
  if (IREE_UNLIKELY(payload.flags != IREE_HAL_REPLAY_SCOPE_FLAG_NONE ||
                    payload.reserved0 != 0 || payload.reserved1 != 0)) {
    return iree_make_status(IREE_STATUS_DATA_LOSS,
                            "replay scope payload reserved fields must be "
                            "zero");
  }
  if (IREE_UNLIKELY(payload.name_length > IREE_HOST_SIZE_MAX ||
                    sizeof(payload) + (iree_host_size_t)payload.name_length !=
                        record->payload.data_length ||
                    payload.name_length == 0)) {
    return iree_make_status(IREE_STATUS_DATA_LOSS,
                            "replay scope payload name length mismatch");
  }
  out_scope->event_type = event_type;
  out_scope->name =
      iree_make_string_view((const char*)record->payload.data + sizeof(payload),
                            (iree_host_size_t)payload.name_length);
  return iree_ok_status();
}

static iree_status_t
iree_hal_replay_plan_prepare_command_buffer_execution_barrier(
    const iree_hal_replay_file_record_t* record,
    iree_hal_replay_plan_command_buffer_execution_barrier_t* out_barrier) {
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_require_payload(
      record, IREE_HAL_REPLAY_PAYLOAD_TYPE_COMMAND_BUFFER_EXECUTION_BARRIER,
      sizeof(iree_hal_replay_command_buffer_execution_barrier_payload_t)));
  iree_hal_replay_command_buffer_execution_barrier_payload_t payload;
  memcpy(&payload, record->payload.data, sizeof(payload));
  iree_host_size_t memory_payloads_size = 0;
  iree_host_size_t buffer_payloads_size = 0;
  iree_host_size_t total_payload_size = 0;
  if (IREE_UNLIKELY(
          payload.memory_barrier_count > IREE_HOST_SIZE_MAX ||
          payload.buffer_barrier_count > IREE_HOST_SIZE_MAX ||
          !iree_host_size_checked_mul(
              (iree_host_size_t)payload.memory_barrier_count,
              sizeof(iree_hal_replay_memory_barrier_payload_t),
              &memory_payloads_size) ||
          !iree_host_size_checked_mul(
              (iree_host_size_t)payload.buffer_barrier_count,
              sizeof(iree_hal_replay_buffer_barrier_payload_t),
              &buffer_payloads_size) ||
          !iree_host_size_checked_add(sizeof(payload), memory_payloads_size,
                                      &total_payload_size) ||
          !iree_host_size_checked_add(total_payload_size, buffer_payloads_size,
                                      &total_payload_size) ||
          total_payload_size != record->payload.data_length)) {
    return iree_make_status(IREE_STATUS_DATA_LOSS,
                            "replay execution barrier payload length mismatch");
  }
  out_barrier->command_buffer_id = record->header.object_id;
  out_barrier->source_stage_mask = payload.source_stage_mask;
  out_barrier->target_stage_mask = payload.target_stage_mask;
  out_barrier->flags = payload.flags;
  out_barrier->memory_barrier_count =
      (iree_host_size_t)payload.memory_barrier_count;
  out_barrier->memory_barrier_payloads =
      (const iree_hal_replay_memory_barrier_payload_t*)(record->payload.data +
                                                        sizeof(payload));
  out_barrier->buffer_barrier_count =
      (iree_host_size_t)payload.buffer_barrier_count;
  out_barrier->buffer_barrier_payloads =
      (const iree_hal_replay_buffer_barrier_payload_t*)(record->payload.data +
                                                        sizeof(payload) +
                                                        memory_payloads_size);
  return iree_ok_status();
}

static iree_status_t iree_hal_replay_plan_prepare_command_buffer_dispatch(
    const iree_hal_replay_file_record_t* record,
    iree_hal_replay_plan_command_buffer_dispatch_t* out_dispatch) {
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_require_payload(
      record, IREE_HAL_REPLAY_PAYLOAD_TYPE_DISPATCH,
      sizeof(iree_hal_replay_dispatch_payload_t)));
  iree_hal_replay_dispatch_payload_t payload;
  memcpy(&payload, record->payload.data, sizeof(payload));
  if (payload.wait_semaphore_count != 0 ||
      payload.signal_semaphore_count != 0) {
    return iree_make_status(IREE_STATUS_UNIMPLEMENTED,
                            "command buffer dispatch semaphore lists are "
                            "reserved for immediate dispatch");
  }
  iree_host_size_t wait_offset = 0;
  iree_host_size_t wait_size = 0;
  iree_host_size_t signal_offset = 0;
  iree_host_size_t signal_size = 0;
  iree_host_size_t constants_offset = 0;
  iree_host_size_t binding_offset = 0;
  iree_host_size_t binding_size = 0;
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_dispatch_layout(
      record, &payload, &wait_offset, &wait_size, &signal_offset, &signal_size,
      &constants_offset, &binding_offset, &binding_size));
  (void)wait_offset;
  (void)wait_size;
  (void)signal_offset;
  (void)signal_size;
  (void)binding_size;
  out_dispatch->command_buffer_id = record->header.object_id;
  out_dispatch->executable_id = payload.executable_id;
  out_dispatch->function_ordinal = payload.function_ordinal;
  memset(&out_dispatch->config, 0, sizeof(out_dispatch->config));
  memcpy(out_dispatch->config.workgroup_size, payload.workgroup_size,
         sizeof(out_dispatch->config.workgroup_size));
  memcpy(out_dispatch->config.workgroup_count, payload.workgroup_count,
         sizeof(out_dispatch->config.workgroup_count));
  out_dispatch->config.dynamic_workgroup_local_memory =
      payload.dynamic_workgroup_local_memory;
  out_dispatch->workgroup_count_ref = payload.workgroup_count_ref;
  out_dispatch->constants =
      iree_make_const_byte_span(record->payload.data + constants_offset,
                                (iree_host_size_t)payload.constants_length);
  out_dispatch->binding_count = (iree_host_size_t)payload.binding_count;
  out_dispatch->binding_payloads =
      (const iree_hal_replay_buffer_ref_payload_t*)(record->payload.data +
                                                    binding_offset);
  out_dispatch->flags = payload.flags;
  return iree_ok_status();
}

static iree_status_t iree_hal_replay_plan_prepare_queue_execute(
    const iree_hal_replay_file_record_t* record,
    iree_hal_replay_plan_queue_execute_t* out_execute) {
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_require_payload(
      record, IREE_HAL_REPLAY_PAYLOAD_TYPE_DEVICE_QUEUE_EXECUTE,
      sizeof(iree_hal_replay_device_queue_execute_payload_t)));
  iree_hal_replay_device_queue_execute_payload_t payload;
  memcpy(&payload, record->payload.data, sizeof(payload));
  iree_host_size_t wait_size = 0;
  iree_host_size_t signal_size = 0;
  iree_host_size_t binding_size = 0;
  iree_host_size_t expected_size = 0;
  if (IREE_UNLIKELY(payload.wait_semaphore_count > IREE_HOST_SIZE_MAX ||
                    payload.signal_semaphore_count > IREE_HOST_SIZE_MAX ||
                    payload.binding_count > IREE_HOST_SIZE_MAX ||
                    !iree_host_size_checked_mul(
                        (iree_host_size_t)payload.wait_semaphore_count,
                        sizeof(iree_hal_replay_semaphore_timepoint_payload_t),
                        &wait_size) ||
                    !iree_host_size_checked_mul(
                        (iree_host_size_t)payload.signal_semaphore_count,
                        sizeof(iree_hal_replay_semaphore_timepoint_payload_t),
                        &signal_size) ||
                    !iree_host_size_checked_mul(
                        (iree_host_size_t)payload.binding_count,
                        sizeof(iree_hal_replay_buffer_ref_payload_t),
                        &binding_size) ||
                    !iree_host_size_checked_add(sizeof(payload), wait_size,
                                                &expected_size) ||
                    !iree_host_size_checked_add(expected_size, signal_size,
                                                &expected_size) ||
                    !iree_host_size_checked_add(expected_size, binding_size,
                                                &expected_size) ||
                    expected_size != record->payload.data_length)) {
    return iree_make_status(IREE_STATUS_DATA_LOSS,
                            "replay queue execute payload length mismatch");
  }
  const uint8_t* cursor = record->payload.data + sizeof(payload);
  out_execute->device_id = record->header.object_id;
  out_execute->command_buffer_id = payload.command_buffer_id;
  out_execute->queue_affinity = payload.queue_affinity;
  out_execute->flags = payload.flags;
  out_execute->wait_semaphore_count =
      (iree_host_size_t)payload.wait_semaphore_count;
  out_execute->wait_semaphore_payloads =
      (const iree_hal_replay_semaphore_timepoint_payload_t*)cursor;
  cursor += wait_size;
  out_execute->signal_semaphore_count =
      (iree_host_size_t)payload.signal_semaphore_count;
  out_execute->signal_semaphore_payloads =
      (const iree_hal_replay_semaphore_timepoint_payload_t*)cursor;
  cursor += signal_size;
  out_execute->binding_count = (iree_host_size_t)payload.binding_count;
  out_execute->binding_payloads =
      (const iree_hal_replay_buffer_ref_payload_t*)cursor;
  return iree_ok_status();
}

static iree_status_t iree_hal_replay_plan_prepare_record(
    iree_hal_replay_plan_record_t* plan_record) {
  const iree_hal_replay_file_record_t* record = &plan_record->file_record;
  plan_record->kind = IREE_HAL_REPLAY_PLAN_RECORD_KIND_GENERIC;
  if (record->header.record_type !=
      IREE_HAL_REPLAY_FILE_RECORD_TYPE_OPERATION) {
    return iree_ok_status();
  }
  if (record->header.status_code != IREE_STATUS_OK) {
    plan_record->kind = IREE_HAL_REPLAY_PLAN_RECORD_KIND_SKIP;
    return iree_ok_status();
  }
  switch (record->header.operation_code) {
    case IREE_HAL_REPLAY_OPERATION_CODE_REPLAY_SCOPE_BEGIN:
      plan_record->kind = IREE_HAL_REPLAY_PLAN_RECORD_KIND_SCOPE;
      return iree_hal_replay_plan_prepare_scope(
          record, IREE_HAL_REPLAY_SCOPE_EVENT_TYPE_BEGIN,
          &plan_record->payload.scope);
    case IREE_HAL_REPLAY_OPERATION_CODE_REPLAY_SCOPE_END:
      plan_record->kind = IREE_HAL_REPLAY_PLAN_RECORD_KIND_SCOPE;
      return iree_hal_replay_plan_prepare_scope(
          record, IREE_HAL_REPLAY_SCOPE_EVENT_TYPE_END,
          &plan_record->payload.scope);
    case IREE_HAL_REPLAY_OPERATION_CODE_COMMAND_BUFFER_EXECUTION_BARRIER:
      plan_record->kind =
          IREE_HAL_REPLAY_PLAN_RECORD_KIND_COMMAND_BUFFER_EXECUTION_BARRIER;
      return iree_hal_replay_plan_prepare_command_buffer_execution_barrier(
          record, &plan_record->payload.command_buffer_execution_barrier);
    case IREE_HAL_REPLAY_OPERATION_CODE_COMMAND_BUFFER_DISPATCH:
      plan_record->kind =
          IREE_HAL_REPLAY_PLAN_RECORD_KIND_COMMAND_BUFFER_DISPATCH;
      return iree_hal_replay_plan_prepare_command_buffer_dispatch(
          record, &plan_record->payload.command_buffer_dispatch);
    case IREE_HAL_REPLAY_OPERATION_CODE_DEVICE_QUEUE_EXECUTE:
      plan_record->kind = IREE_HAL_REPLAY_PLAN_RECORD_KIND_QUEUE_EXECUTE;
      return iree_hal_replay_plan_prepare_queue_execute(
          record, &plan_record->payload.queue_execute);
    default:
      return iree_ok_status();
  }
}

IREE_API_EXPORT void iree_hal_replay_plan_destroy(
    iree_hal_replay_plan_t* plan) {
  if (!plan) return;
  iree_allocator_t host_allocator = plan->host_allocator;
  iree_allocator_free(host_allocator, plan->records);
  iree_allocator_free(host_allocator, plan);
}

IREE_API_EXPORT iree_status_t iree_hal_replay_plan_create(
    iree_const_byte_span_t file_contents, iree_allocator_t host_allocator,
    iree_hal_replay_plan_t** out_plan) {
  IREE_ASSERT_ARGUMENT(out_plan);
  *out_plan = NULL;

  iree_hal_replay_file_header_t file_header;
  iree_host_size_t first_record_offset = 0;
  IREE_RETURN_IF_ERROR(iree_hal_replay_file_parse_header(
      file_contents, &file_header, &first_record_offset));
  iree_const_byte_span_t valid_contents = file_contents;
  if (file_header.file_length != 0) {
    valid_contents.data_length = (iree_host_size_t)file_header.file_length;
  }

  iree_host_size_t record_count = 0;
  iree_hal_replay_object_id_t max_object_id = IREE_HAL_REPLAY_OBJECT_ID_NONE;
  iree_host_size_t offset = first_record_offset;
  uint64_t expected_sequence_ordinal = 0;
  while (offset < valid_contents.data_length) {
    iree_hal_replay_file_record_t record;
    IREE_RETURN_IF_ERROR(iree_hal_replay_file_parse_record(
        valid_contents, offset, &record, &offset));
    if (IREE_UNLIKELY(record.header.sequence_ordinal !=
                      expected_sequence_ordinal++)) {
      return iree_make_status(IREE_STATUS_DATA_LOSS,
                              "replay record sequence ordinal mismatch");
    }
    if (record.header.object_id > max_object_id) {
      max_object_id = record.header.object_id;
    }
    if (record.header.related_object_id > max_object_id) {
      max_object_id = record.header.related_object_id;
    }
    ++record_count;
  }
  if (IREE_UNLIKELY(max_object_id >= IREE_HOST_SIZE_MAX)) {
    return iree_make_status(IREE_STATUS_OUT_OF_RANGE,
                            "replay object id exceeds host size");
  }

  iree_hal_replay_plan_t* plan = NULL;
  IREE_RETURN_IF_ERROR(
      iree_allocator_malloc(host_allocator, sizeof(*plan), (void**)&plan));
  memset(plan, 0, sizeof(*plan));
  plan->file_contents = valid_contents;
  plan->host_allocator = host_allocator;
  plan->object_capacity = (iree_host_size_t)max_object_id + 1;
  plan->record_count = record_count;

  iree_status_t status = iree_ok_status();
  if (record_count != 0) {
    iree_host_size_t records_size = 0;
    if (IREE_UNLIKELY(!iree_host_size_checked_mul(
            record_count, sizeof(*plan->records), &records_size))) {
      status = iree_make_status(IREE_STATUS_OUT_OF_RANGE,
                                "replay plan record table size overflow");
    }
    if (iree_status_is_ok(status)) {
      status = iree_allocator_malloc(host_allocator, records_size,
                                     (void**)&plan->records);
    }
    if (iree_status_is_ok(status)) {
      memset(plan->records, 0, records_size);
    }
  }

  offset = first_record_offset;
  for (iree_host_size_t i = 0; i < record_count && iree_status_is_ok(status);
       ++i) {
    iree_hal_replay_plan_record_t* plan_record = &plan->records[i];
    status = iree_hal_replay_file_parse_record(
        valid_contents, offset, &plan_record->file_record, &offset);
    if (iree_status_is_ok(status)) {
      status = iree_hal_replay_plan_prepare_record(plan_record);
    }
  }
  if (iree_status_is_ok(status)) {
    *out_plan = plan;
  } else {
    iree_hal_replay_plan_destroy(plan);
  }
  return status;
}

static iree_status_t iree_hal_replay_executor_command_buffer_dispatch(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_file_record_t* record) {
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_require_payload(
      record, IREE_HAL_REPLAY_PAYLOAD_TYPE_DISPATCH,
      sizeof(iree_hal_replay_dispatch_payload_t)));
  iree_hal_replay_dispatch_payload_t payload;
  memcpy(&payload, record->payload.data, sizeof(payload));
  if (payload.wait_semaphore_count != 0 ||
      payload.signal_semaphore_count != 0) {
    return iree_make_status(IREE_STATUS_UNIMPLEMENTED,
                            "command buffer dispatch semaphore lists are "
                            "reserved for immediate dispatch");
  }
  iree_host_size_t wait_offset = 0;
  iree_host_size_t wait_size = 0;
  iree_host_size_t signal_offset = 0;
  iree_host_size_t signal_size = 0;
  iree_host_size_t constants_offset = 0;
  iree_host_size_t binding_offset = 0;
  iree_host_size_t binding_size = 0;
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_dispatch_layout(
      record, &payload, &wait_offset, &wait_size, &signal_offset, &signal_size,
      &constants_offset, &binding_offset, &binding_size));
  iree_const_byte_span_t constants =
      iree_make_const_byte_span(record->payload.data + constants_offset,
                                (iree_host_size_t)payload.constants_length);
  const iree_hal_replay_buffer_ref_payload_t* binding_payloads =
      (const iree_hal_replay_buffer_ref_payload_t*)(record->payload.data +
                                                    binding_offset);
  iree_hal_replay_buffer_ref_list_storage_t binding_storage = {0};
  IREE_RETURN_IF_ERROR(iree_hal_replay_buffer_ref_list_storage_initialize(
      executor, (iree_host_size_t)payload.binding_count, &binding_storage));
  iree_status_t status = iree_ok_status();
  for (iree_host_size_t i = 0;
       i < payload.binding_count && iree_status_is_ok(status); ++i) {
    status = iree_hal_replay_executor_make_buffer_ref(
        executor, &binding_payloads[i], &binding_storage.values[i]);
  }
  iree_hal_replay_object_entry_t* command_buffer_entry = NULL;
  if (iree_status_is_ok(status)) {
    status = iree_hal_replay_executor_lookup(
        executor, record->header.object_id,
        IREE_HAL_REPLAY_OBJECT_TYPE_COMMAND_BUFFER, &command_buffer_entry);
  }
  iree_hal_replay_object_entry_t* executable_entry = NULL;
  if (iree_status_is_ok(status)) {
    status = iree_hal_replay_executor_lookup(
        executor, payload.executable_id, IREE_HAL_REPLAY_OBJECT_TYPE_EXECUTABLE,
        &executable_entry);
  }
  if (iree_status_is_ok(status)) {
    iree_hal_dispatch_config_t config;
    memset(&config, 0, sizeof(config));
    memcpy(config.workgroup_size, payload.workgroup_size,
           sizeof(config.workgroup_size));
    memcpy(config.workgroup_count, payload.workgroup_count,
           sizeof(config.workgroup_count));
    config.dynamic_workgroup_local_memory =
        payload.dynamic_workgroup_local_memory;
    iree_hal_executable_function_t function =
        iree_hal_executable_function_invalid();
    status = iree_hal_replay_executor_make_buffer_ref(
        executor, &payload.workgroup_count_ref, &config.workgroup_count_ref);
    if (iree_status_is_ok(status)) {
      status = iree_hal_replay_executor_resolve_function(
          payload.executable_id, &executable_entry->value.executable,
          payload.function_ordinal, &function);
    }
    if (iree_status_is_ok(status)) {
      status = iree_hal_command_buffer_dispatch(
          command_buffer_entry->value.command_buffer,
          executable_entry->value.executable.handle, function, config,
          constants, binding_storage.list, payload.flags);
    }
  }
  iree_hal_replay_buffer_ref_list_storage_deinitialize(
      &binding_storage, executor->host_allocator);
  return status;
}

static iree_status_t iree_hal_replay_executor_queue_dispatch(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_file_record_t* record) {
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_require_payload(
      record, IREE_HAL_REPLAY_PAYLOAD_TYPE_DISPATCH,
      sizeof(iree_hal_replay_dispatch_payload_t)));
  iree_hal_replay_dispatch_payload_t payload;
  memcpy(&payload, record->payload.data, sizeof(payload));
  iree_host_size_t wait_offset = 0;
  iree_host_size_t wait_size = 0;
  iree_host_size_t signal_offset = 0;
  iree_host_size_t signal_size = 0;
  iree_host_size_t constants_offset = 0;
  iree_host_size_t binding_offset = 0;
  iree_host_size_t binding_size = 0;
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_dispatch_layout(
      record, &payload, &wait_offset, &wait_size, &signal_offset, &signal_size,
      &constants_offset, &binding_offset, &binding_size));
  iree_const_byte_span_t constants =
      iree_make_const_byte_span(record->payload.data + constants_offset,
                                (iree_host_size_t)payload.constants_length);
  const iree_hal_replay_buffer_ref_payload_t* binding_payloads =
      (const iree_hal_replay_buffer_ref_payload_t*)(record->payload.data +
                                                    binding_offset);

  iree_hal_replay_semaphore_list_storage_t wait_storage;
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_make_semaphore_list(
      executor,
      iree_make_const_byte_span(record->payload.data + wait_offset, wait_size),
      (iree_host_size_t)payload.wait_semaphore_count, &wait_storage));
  iree_hal_replay_semaphore_list_storage_t signal_storage;
  iree_status_t status = iree_hal_replay_executor_make_semaphore_list(
      executor,
      iree_make_const_byte_span(record->payload.data + signal_offset,
                                signal_size),
      (iree_host_size_t)payload.signal_semaphore_count, &signal_storage);

  iree_hal_replay_buffer_ref_list_storage_t binding_storage = {0};
  if (iree_status_is_ok(status)) {
    status = iree_hal_replay_buffer_ref_list_storage_initialize(
        executor, (iree_host_size_t)payload.binding_count, &binding_storage);
  }
  for (iree_host_size_t i = 0;
       i < payload.binding_count && iree_status_is_ok(status); ++i) {
    status = iree_hal_replay_executor_make_buffer_ref(
        executor, &binding_payloads[i], &binding_storage.values[i]);
  }
  iree_hal_replay_object_entry_t* device_entry = NULL;
  if (iree_status_is_ok(status)) {
    status = iree_hal_replay_executor_lookup(executor, record->header.object_id,
                                             IREE_HAL_REPLAY_OBJECT_TYPE_DEVICE,
                                             &device_entry);
  }
  iree_hal_replay_object_entry_t* executable_entry = NULL;
  if (iree_status_is_ok(status)) {
    status = iree_hal_replay_executor_lookup(
        executor, payload.executable_id, IREE_HAL_REPLAY_OBJECT_TYPE_EXECUTABLE,
        &executable_entry);
  }
  if (iree_status_is_ok(status)) {
    iree_hal_dispatch_config_t config;
    memset(&config, 0, sizeof(config));
    memcpy(config.workgroup_size, payload.workgroup_size,
           sizeof(config.workgroup_size));
    memcpy(config.workgroup_count, payload.workgroup_count,
           sizeof(config.workgroup_count));
    config.dynamic_workgroup_local_memory =
        payload.dynamic_workgroup_local_memory;
    iree_hal_executable_function_t function =
        iree_hal_executable_function_invalid();
    status = iree_hal_replay_executor_make_buffer_ref(
        executor, &payload.workgroup_count_ref, &config.workgroup_count_ref);
    if (iree_status_is_ok(status)) {
      status = iree_hal_replay_executor_resolve_function(
          payload.executable_id, &executable_entry->value.executable,
          payload.function_ordinal, &function);
    }
    if (iree_status_is_ok(status)) {
      status = iree_hal_device_queue_dispatch(
          device_entry->value.device, payload.queue_affinity, wait_storage.list,
          signal_storage.list, executable_entry->value.executable.handle,
          function, config, constants, binding_storage.list, payload.flags);
    }
  }
  if (iree_status_is_ok(status)) {
    status = iree_hal_replay_executor_flush_and_wait(device_entry->value.device,
                                                     payload.queue_affinity,
                                                     signal_storage.list);
  }

  iree_hal_replay_buffer_ref_list_storage_deinitialize(
      &binding_storage, executor->host_allocator);
  iree_hal_replay_semaphore_list_storage_deinitialize(&signal_storage,
                                                      executor->host_allocator);
  iree_hal_replay_semaphore_list_storage_deinitialize(&wait_storage,
                                                      executor->host_allocator);
  return status;
}

static iree_status_t iree_hal_replay_executor_command_buffer_fill_buffer(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_file_record_t* record) {
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_require_payload(
      record, IREE_HAL_REPLAY_PAYLOAD_TYPE_COMMAND_BUFFER_FILL_BUFFER,
      sizeof(iree_hal_replay_command_buffer_fill_buffer_payload_t)));
  iree_hal_replay_command_buffer_fill_buffer_payload_t payload;
  memcpy(&payload, record->payload.data, sizeof(payload));
  if (IREE_UNLIKELY(payload.pattern_length > IREE_HOST_SIZE_MAX ||
                    sizeof(payload) +
                            (iree_host_size_t)payload.pattern_length !=
                        record->payload.data_length)) {
    return iree_make_status(IREE_STATUS_DATA_LOSS,
                            "replay command buffer fill payload length "
                            "mismatch");
  }
  iree_hal_replay_object_entry_t* command_buffer_entry = NULL;
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_lookup(
      executor, record->header.object_id,
      IREE_HAL_REPLAY_OBJECT_TYPE_COMMAND_BUFFER, &command_buffer_entry));
  iree_hal_buffer_ref_t target_ref;
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_make_buffer_ref(
      executor, &payload.target_ref, &target_ref));
  iree_const_byte_span_t pattern =
      iree_make_const_byte_span(record->payload.data + sizeof(payload),
                                (iree_host_size_t)payload.pattern_length);
  return iree_hal_command_buffer_fill_buffer(
      command_buffer_entry->value.command_buffer, target_ref, pattern.data,
      pattern.data_length, payload.flags);
}

static iree_status_t iree_hal_replay_executor_command_buffer_update_buffer(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_file_record_t* record) {
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_require_payload(
      record, IREE_HAL_REPLAY_PAYLOAD_TYPE_COMMAND_BUFFER_UPDATE_BUFFER,
      sizeof(iree_hal_replay_command_buffer_update_buffer_payload_t)));
  iree_hal_replay_command_buffer_update_buffer_payload_t payload;
  memcpy(&payload, record->payload.data, sizeof(payload));
  if (IREE_UNLIKELY(payload.data_length > IREE_HOST_SIZE_MAX ||
                    sizeof(payload) + (iree_host_size_t)payload.data_length !=
                        record->payload.data_length)) {
    return iree_make_status(IREE_STATUS_DATA_LOSS,
                            "replay command buffer update payload length "
                            "mismatch");
  }
  iree_hal_replay_object_entry_t* command_buffer_entry = NULL;
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_lookup(
      executor, record->header.object_id,
      IREE_HAL_REPLAY_OBJECT_TYPE_COMMAND_BUFFER, &command_buffer_entry));
  iree_hal_buffer_ref_t target_ref;
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_make_buffer_ref(
      executor, &payload.target_ref, &target_ref));
  if (payload.data_length != target_ref.length) {
    return iree_make_status(IREE_STATUS_DATA_LOSS,
                            "replay command buffer update data length does not "
                            "match target length");
  }
  iree_const_byte_span_t data =
      iree_make_const_byte_span(record->payload.data + sizeof(payload),
                                (iree_host_size_t)payload.data_length);
  return iree_hal_command_buffer_update_buffer(
      command_buffer_entry->value.command_buffer, data.data,
      /*source_offset=*/0, target_ref, payload.flags);
}

static iree_status_t iree_hal_replay_executor_command_buffer_copy_buffer(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_file_record_t* record) {
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_require_payload(
      record, IREE_HAL_REPLAY_PAYLOAD_TYPE_COMMAND_BUFFER_COPY_BUFFER,
      sizeof(iree_hal_replay_command_buffer_copy_buffer_payload_t)));
  iree_hal_replay_command_buffer_copy_buffer_payload_t payload;
  memcpy(&payload, record->payload.data, sizeof(payload));
  iree_hal_replay_object_entry_t* command_buffer_entry = NULL;
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_lookup(
      executor, record->header.object_id,
      IREE_HAL_REPLAY_OBJECT_TYPE_COMMAND_BUFFER, &command_buffer_entry));
  iree_hal_buffer_ref_t source_ref;
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_make_buffer_ref(
      executor, &payload.source_ref, &source_ref));
  iree_hal_buffer_ref_t target_ref;
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_make_buffer_ref(
      executor, &payload.target_ref, &target_ref));
  return iree_hal_command_buffer_copy_buffer(
      command_buffer_entry->value.command_buffer, source_ref, target_ref,
      payload.flags);
}

static iree_status_t iree_hal_replay_executor_queue_execute(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_file_record_t* record) {
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_require_payload(
      record, IREE_HAL_REPLAY_PAYLOAD_TYPE_DEVICE_QUEUE_EXECUTE,
      sizeof(iree_hal_replay_device_queue_execute_payload_t)));
  iree_hal_replay_device_queue_execute_payload_t payload;
  memcpy(&payload, record->payload.data, sizeof(payload));
  iree_host_size_t wait_size = 0;
  iree_host_size_t signal_size = 0;
  iree_host_size_t binding_size = 0;
  iree_host_size_t expected_size = 0;
  if (IREE_UNLIKELY(payload.wait_semaphore_count > IREE_HOST_SIZE_MAX ||
                    payload.signal_semaphore_count > IREE_HOST_SIZE_MAX ||
                    payload.binding_count > IREE_HOST_SIZE_MAX ||
                    !iree_host_size_checked_mul(
                        (iree_host_size_t)payload.wait_semaphore_count,
                        sizeof(iree_hal_replay_semaphore_timepoint_payload_t),
                        &wait_size) ||
                    !iree_host_size_checked_mul(
                        (iree_host_size_t)payload.signal_semaphore_count,
                        sizeof(iree_hal_replay_semaphore_timepoint_payload_t),
                        &signal_size) ||
                    !iree_host_size_checked_mul(
                        (iree_host_size_t)payload.binding_count,
                        sizeof(iree_hal_replay_buffer_ref_payload_t),
                        &binding_size) ||
                    !iree_host_size_checked_add(sizeof(payload), wait_size,
                                                &expected_size) ||
                    !iree_host_size_checked_add(expected_size, signal_size,
                                                &expected_size) ||
                    !iree_host_size_checked_add(expected_size, binding_size,
                                                &expected_size) ||
                    expected_size != record->payload.data_length)) {
    return iree_make_status(IREE_STATUS_DATA_LOSS,
                            "replay queue execute payload length mismatch");
  }
  const uint8_t* cursor = record->payload.data + sizeof(payload);
  iree_hal_replay_semaphore_list_storage_t wait_storage;
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_make_semaphore_list(
      executor, iree_make_const_byte_span(cursor, wait_size),
      (iree_host_size_t)payload.wait_semaphore_count, &wait_storage));
  cursor += wait_size;
  iree_hal_replay_semaphore_list_storage_t signal_storage;
  iree_status_t status = iree_hal_replay_executor_make_semaphore_list(
      executor, iree_make_const_byte_span(cursor, signal_size),
      (iree_host_size_t)payload.signal_semaphore_count, &signal_storage);
  cursor += signal_size;

  iree_hal_replay_buffer_binding_table_storage_t binding_storage = {0};
  if (iree_status_is_ok(status)) {
    status = iree_hal_replay_buffer_binding_table_storage_initialize(
        executor, (iree_host_size_t)payload.binding_count, &binding_storage);
  }
  const iree_hal_replay_buffer_ref_payload_t* binding_payloads =
      (const iree_hal_replay_buffer_ref_payload_t*)cursor;
  for (iree_host_size_t i = 0;
       i < payload.binding_count && iree_status_is_ok(status); ++i) {
    iree_hal_buffer_ref_t ref;
    status = iree_hal_replay_executor_make_buffer_ref(
        executor, &binding_payloads[i], &ref);
    if (iree_status_is_ok(status)) {
      binding_storage.bindings[i] = (iree_hal_buffer_binding_t){
          .buffer = ref.buffer,
          .offset = ref.offset,
          .length = ref.length,
      };
    }
  }

  iree_hal_replay_object_entry_t* device_entry = NULL;
  if (iree_status_is_ok(status)) {
    status = iree_hal_replay_executor_lookup(executor, record->header.object_id,
                                             IREE_HAL_REPLAY_OBJECT_TYPE_DEVICE,
                                             &device_entry);
  }
  iree_hal_replay_object_entry_t* command_buffer_entry = NULL;
  if (iree_status_is_ok(status) &&
      payload.command_buffer_id != IREE_HAL_REPLAY_OBJECT_ID_NONE) {
    status = iree_hal_replay_executor_lookup(
        executor, payload.command_buffer_id,
        IREE_HAL_REPLAY_OBJECT_TYPE_COMMAND_BUFFER, &command_buffer_entry);
  }
  if (iree_status_is_ok(status)) {
    if (command_buffer_entry) {
      status = iree_hal_device_queue_execute(
          device_entry->value.device, payload.queue_affinity, wait_storage.list,
          signal_storage.list, command_buffer_entry->value.command_buffer,
          binding_storage.table, payload.flags);
    } else if (payload.binding_count == 0) {
      status = iree_hal_device_queue_barrier(
          device_entry->value.device, payload.queue_affinity, wait_storage.list,
          signal_storage.list, payload.flags);
    } else {
      status = iree_make_status(
          IREE_STATUS_DATA_LOSS,
          "replay queue barrier payload unexpectedly has bindings");
    }
  }
  if (iree_status_is_ok(status)) {
    status = iree_hal_device_queue_flush(device_entry->value.device,
                                         payload.queue_affinity);
  }
  if (iree_status_is_ok(status) && signal_storage.list.count != 0) {
    status = iree_hal_semaphore_list_wait(signal_storage.list,
                                          iree_infinite_timeout(),
                                          IREE_ASYNC_WAIT_FLAG_NONE);
  }

  iree_hal_replay_buffer_binding_table_storage_deinitialize(
      &binding_storage, executor->host_allocator);
  iree_hal_replay_semaphore_list_storage_deinitialize(&signal_storage,
                                                      executor->host_allocator);
  iree_hal_replay_semaphore_list_storage_deinitialize(&wait_storage,
                                                      executor->host_allocator);
  return status;
}

static iree_status_t iree_hal_replay_plan_execute_scope(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_plan_record_t* record) {
  iree_hal_replay_scope_event_callback_t callback =
      executor->options->scope_event_callback;
  if (!callback.fn) return iree_ok_status();
  const iree_hal_replay_plan_scope_t* scope = &record->payload.scope;
  iree_hal_replay_scope_event_t event = {
      .sequence_ordinal = record->file_record.header.sequence_ordinal,
      .type = scope->event_type,
      .name = scope->name,
  };
  return callback.fn(callback.user_data, &event);
}

static iree_status_t
iree_hal_replay_plan_execute_command_buffer_execution_barrier(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_plan_command_buffer_execution_barrier_t* barrier) {
  iree_hal_memory_barrier_t inline_memory_barriers
      [IREE_HAL_REPLAY_INLINE_MEMORY_BARRIER_LIST_CAPACITY];
  iree_hal_memory_barrier_t* memory_barriers = NULL;
  bool memory_barriers_allocated = false;
  if (barrier->memory_barrier_count <=
      IREE_HAL_REPLAY_INLINE_MEMORY_BARRIER_LIST_CAPACITY) {
    memory_barriers = inline_memory_barriers;
  } else {
    iree_host_size_t memory_barriers_size = 0;
    if (IREE_UNLIKELY(!iree_host_size_checked_mul(barrier->memory_barrier_count,
                                                  sizeof(*memory_barriers),
                                                  &memory_barriers_size))) {
      return iree_make_status(
          IREE_STATUS_OUT_OF_RANGE,
          "replay execution barrier memory barrier count overflow");
    }
    IREE_RETURN_IF_ERROR(iree_allocator_malloc(executor->host_allocator,
                                               memory_barriers_size,
                                               (void**)&memory_barriers));
    memory_barriers_allocated = true;
  }
  for (iree_host_size_t i = 0; i < barrier->memory_barrier_count; ++i) {
    memory_barriers[i].source_scope =
        barrier->memory_barrier_payloads[i].source_scope;
    memory_barriers[i].target_scope =
        barrier->memory_barrier_payloads[i].target_scope;
  }

  iree_hal_buffer_barrier_t inline_buffer_barriers
      [IREE_HAL_REPLAY_INLINE_BUFFER_BARRIER_LIST_CAPACITY];
  iree_hal_buffer_barrier_t* buffer_barriers = NULL;
  iree_status_t status = iree_ok_status();
  bool buffer_barriers_allocated = false;
  if (barrier->buffer_barrier_count <=
      IREE_HAL_REPLAY_INLINE_BUFFER_BARRIER_LIST_CAPACITY) {
    buffer_barriers = inline_buffer_barriers;
  } else {
    iree_host_size_t buffer_barriers_size = 0;
    if (IREE_UNLIKELY(!iree_host_size_checked_mul(barrier->buffer_barrier_count,
                                                  sizeof(*buffer_barriers),
                                                  &buffer_barriers_size))) {
      status = iree_make_status(
          IREE_STATUS_OUT_OF_RANGE,
          "replay execution barrier buffer barrier count overflow");
    }
    if (iree_status_is_ok(status)) {
      status =
          iree_allocator_malloc(executor->host_allocator, buffer_barriers_size,
                                (void**)&buffer_barriers);
    }
    buffer_barriers_allocated = iree_status_is_ok(status);
  }
  for (iree_host_size_t i = 0;
       i < barrier->buffer_barrier_count && iree_status_is_ok(status); ++i) {
    buffer_barriers[i].source_scope =
        barrier->buffer_barrier_payloads[i].source_scope;
    buffer_barriers[i].target_scope =
        barrier->buffer_barrier_payloads[i].target_scope;
    status = iree_hal_replay_executor_make_buffer_ref(
        executor, &barrier->buffer_barrier_payloads[i].buffer_ref,
        &buffer_barriers[i].buffer_ref);
  }

  iree_hal_replay_object_entry_t* command_buffer_entry = NULL;
  if (iree_status_is_ok(status)) {
    status = iree_hal_replay_executor_lookup(
        executor, barrier->command_buffer_id,
        IREE_HAL_REPLAY_OBJECT_TYPE_COMMAND_BUFFER, &command_buffer_entry);
  }
  if (iree_status_is_ok(status)) {
    status = iree_hal_command_buffer_execution_barrier(
        command_buffer_entry->value.command_buffer, barrier->source_stage_mask,
        barrier->target_stage_mask, barrier->flags,
        barrier->memory_barrier_count, memory_barriers,
        barrier->buffer_barrier_count, buffer_barriers);
  }
  if (buffer_barriers_allocated) {
    iree_allocator_free(executor->host_allocator, buffer_barriers);
  }
  if (memory_barriers_allocated) {
    iree_allocator_free(executor->host_allocator, memory_barriers);
  }
  return status;
}

static iree_status_t iree_hal_replay_plan_execute_command_buffer_dispatch(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_plan_command_buffer_dispatch_t* dispatch) {
  iree_hal_replay_buffer_ref_list_storage_t binding_storage = {0};
  IREE_RETURN_IF_ERROR(iree_hal_replay_buffer_ref_list_storage_initialize(
      executor, dispatch->binding_count, &binding_storage));
  iree_status_t status = iree_ok_status();
  for (iree_host_size_t i = 0;
       i < dispatch->binding_count && iree_status_is_ok(status); ++i) {
    status = iree_hal_replay_executor_make_buffer_ref(
        executor, &dispatch->binding_payloads[i], &binding_storage.values[i]);
  }

  iree_hal_replay_object_entry_t* command_buffer_entry = NULL;
  if (iree_status_is_ok(status)) {
    status = iree_hal_replay_executor_lookup(
        executor, dispatch->command_buffer_id,
        IREE_HAL_REPLAY_OBJECT_TYPE_COMMAND_BUFFER, &command_buffer_entry);
  }
  iree_hal_replay_object_entry_t* executable_entry = NULL;
  if (iree_status_is_ok(status)) {
    status = iree_hal_replay_executor_lookup(
        executor, dispatch->executable_id,
        IREE_HAL_REPLAY_OBJECT_TYPE_EXECUTABLE, &executable_entry);
  }
  if (iree_status_is_ok(status)) {
    iree_hal_dispatch_config_t config = dispatch->config;
    iree_hal_executable_function_t function =
        iree_hal_executable_function_invalid();
    status = iree_hal_replay_executor_make_buffer_ref(
        executor, &dispatch->workgroup_count_ref, &config.workgroup_count_ref);
    if (iree_status_is_ok(status)) {
      status = iree_hal_replay_executor_resolve_function(
          dispatch->executable_id, &executable_entry->value.executable,
          dispatch->function_ordinal, &function);
    }
    if (iree_status_is_ok(status)) {
      status = iree_hal_command_buffer_dispatch(
          command_buffer_entry->value.command_buffer,
          executable_entry->value.executable.handle, function, config,
          dispatch->constants, binding_storage.list, dispatch->flags);
    }
  }
  iree_hal_replay_buffer_ref_list_storage_deinitialize(
      &binding_storage, executor->host_allocator);
  return status;
}

static iree_status_t iree_hal_replay_plan_execute_queue_execute(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_plan_queue_execute_t* queue_execute) {
  iree_hal_replay_semaphore_list_storage_t wait_storage;
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_make_semaphore_list(
      executor,
      iree_make_const_byte_span(
          (const uint8_t*)queue_execute->wait_semaphore_payloads,
          queue_execute->wait_semaphore_count *
              sizeof(*queue_execute->wait_semaphore_payloads)),
      queue_execute->wait_semaphore_count, &wait_storage));
  iree_hal_replay_semaphore_list_storage_t signal_storage;
  iree_status_t status = iree_hal_replay_executor_make_semaphore_list(
      executor,
      iree_make_const_byte_span(
          (const uint8_t*)queue_execute->signal_semaphore_payloads,
          queue_execute->signal_semaphore_count *
              sizeof(*queue_execute->signal_semaphore_payloads)),
      queue_execute->signal_semaphore_count, &signal_storage);

  iree_hal_replay_buffer_binding_table_storage_t binding_storage = {0};
  if (iree_status_is_ok(status)) {
    status = iree_hal_replay_buffer_binding_table_storage_initialize(
        executor, queue_execute->binding_count, &binding_storage);
  }
  for (iree_host_size_t i = 0;
       i < queue_execute->binding_count && iree_status_is_ok(status); ++i) {
    iree_hal_buffer_ref_t ref;
    status = iree_hal_replay_executor_make_buffer_ref(
        executor, &queue_execute->binding_payloads[i], &ref);
    if (iree_status_is_ok(status)) {
      binding_storage.bindings[i] = (iree_hal_buffer_binding_t){
          .buffer = ref.buffer,
          .offset = ref.offset,
          .length = ref.length,
      };
    }
  }

  iree_hal_replay_object_entry_t* device_entry = NULL;
  if (iree_status_is_ok(status)) {
    status = iree_hal_replay_executor_lookup(executor, queue_execute->device_id,
                                             IREE_HAL_REPLAY_OBJECT_TYPE_DEVICE,
                                             &device_entry);
  }
  iree_hal_replay_object_entry_t* command_buffer_entry = NULL;
  if (iree_status_is_ok(status) &&
      queue_execute->command_buffer_id != IREE_HAL_REPLAY_OBJECT_ID_NONE) {
    status = iree_hal_replay_executor_lookup(
        executor, queue_execute->command_buffer_id,
        IREE_HAL_REPLAY_OBJECT_TYPE_COMMAND_BUFFER, &command_buffer_entry);
  }
  if (iree_status_is_ok(status)) {
    if (command_buffer_entry) {
      status = iree_hal_device_queue_execute(
          device_entry->value.device, queue_execute->queue_affinity,
          wait_storage.list, signal_storage.list,
          command_buffer_entry->value.command_buffer, binding_storage.table,
          queue_execute->flags);
    } else if (queue_execute->binding_count == 0) {
      status = iree_hal_device_queue_barrier(
          device_entry->value.device, queue_execute->queue_affinity,
          wait_storage.list, signal_storage.list, queue_execute->flags);
    } else {
      status = iree_make_status(
          IREE_STATUS_DATA_LOSS,
          "replay queue barrier payload unexpectedly has bindings");
    }
  }
  if (iree_status_is_ok(status)) {
    status = iree_hal_device_queue_flush(device_entry->value.device,
                                         queue_execute->queue_affinity);
  }
  if (iree_status_is_ok(status) && signal_storage.list.count != 0) {
    status = iree_hal_semaphore_list_wait(signal_storage.list,
                                          iree_infinite_timeout(),
                                          IREE_ASYNC_WAIT_FLAG_NONE);
  }

  iree_hal_replay_buffer_binding_table_storage_deinitialize(
      &binding_storage, executor->host_allocator);
  iree_hal_replay_semaphore_list_storage_deinitialize(&signal_storage,
                                                      executor->host_allocator);
  iree_hal_replay_semaphore_list_storage_deinitialize(&wait_storage,
                                                      executor->host_allocator);
  return status;
}

static iree_status_t iree_hal_replay_executor_replay_operation(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_file_record_t* record) {
  if (IREE_UNLIKELY(record->header.status_code != IREE_STATUS_OK)) {
    // Failed calls produced no replay object and are preserved only to explain
    // host fallback branches. Replay follows the later successful calls the
    // original host issued after observing the failure.
    return iree_ok_status();
  }
  switch (record->header.operation_code) {
    case IREE_HAL_REPLAY_OPERATION_CODE_REPLAY_SCOPE_BEGIN:
      return iree_hal_replay_executor_scope_event(
          executor, record, IREE_HAL_REPLAY_SCOPE_EVENT_TYPE_BEGIN);
    case IREE_HAL_REPLAY_OPERATION_CODE_REPLAY_SCOPE_END:
      return iree_hal_replay_executor_scope_event(
          executor, record, IREE_HAL_REPLAY_SCOPE_EVENT_TYPE_END);
    case IREE_HAL_REPLAY_OPERATION_CODE_DEVICE_TRIM:
    case IREE_HAL_REPLAY_OPERATION_CODE_DEVICE_ASSIGN_TOPOLOGY_INFO:
    case IREE_HAL_REPLAY_OPERATION_CODE_DEVICE_QUERY_I64:
    case IREE_HAL_REPLAY_OPERATION_CODE_DEVICE_QUERY_CAPABILITIES:
    case IREE_HAL_REPLAY_OPERATION_CODE_DEVICE_QUERY_QUEUE_POOL_BACKEND:
    case IREE_HAL_REPLAY_OPERATION_CODE_DEVICE_PROFILING_BEGIN:
    case IREE_HAL_REPLAY_OPERATION_CODE_DEVICE_PROFILING_FLUSH:
    case IREE_HAL_REPLAY_OPERATION_CODE_DEVICE_PROFILING_END:
    case IREE_HAL_REPLAY_OPERATION_CODE_DEVICE_EXTERNAL_CAPTURE_BEGIN:
    case IREE_HAL_REPLAY_OPERATION_CODE_DEVICE_EXTERNAL_CAPTURE_END:
    case IREE_HAL_REPLAY_OPERATION_CODE_ALLOCATOR_TRIM:
    case IREE_HAL_REPLAY_OPERATION_CODE_ALLOCATOR_QUERY_MEMORY_HEAPS:
    case IREE_HAL_REPLAY_OPERATION_CODE_BUFFER_INVALIDATE_RANGE:
    case IREE_HAL_REPLAY_OPERATION_CODE_COMMAND_BUFFER_BEGIN_DEBUG_GROUP:
    case IREE_HAL_REPLAY_OPERATION_CODE_COMMAND_BUFFER_END_DEBUG_GROUP:
    case IREE_HAL_REPLAY_OPERATION_CODE_COMMAND_BUFFER_ADVISE_BUFFER:
    case IREE_HAL_REPLAY_OPERATION_CODE_EXECUTABLE_FUNCTION_COUNT:
    case IREE_HAL_REPLAY_OPERATION_CODE_EXECUTABLE_FUNCTION_INFO:
    case IREE_HAL_REPLAY_OPERATION_CODE_EXECUTABLE_FUNCTION_PARAMETERS:
    case IREE_HAL_REPLAY_OPERATION_CODE_EXECUTABLE_LOOKUP_FUNCTION_BY_NAME:
    case IREE_HAL_REPLAY_OPERATION_CODE_BUFFER_MAP_RANGE:
      return iree_ok_status();
    case IREE_HAL_REPLAY_OPERATION_CODE_DEVICE_CREATE_EXECUTABLE_CACHE:
      return iree_hal_replay_executor_create_executable_cache(executor, record);
    case IREE_HAL_REPLAY_OPERATION_CODE_EXECUTABLE_CACHE_PREPARE_EXECUTABLE:
      return iree_hal_replay_executor_prepare_executable(executor, record);
    case IREE_HAL_REPLAY_OPERATION_CODE_DEVICE_CREATE_COMMAND_BUFFER:
      return iree_hal_replay_executor_create_command_buffer(executor, record);
    case IREE_HAL_REPLAY_OPERATION_CODE_DEVICE_CREATE_EVENT:
      return iree_hal_replay_executor_create_event(executor, record);
    case IREE_HAL_REPLAY_OPERATION_CODE_DEVICE_IMPORT_FILE:
      return iree_hal_replay_executor_import_file(executor, record);
    case IREE_HAL_REPLAY_OPERATION_CODE_DEVICE_CREATE_SEMAPHORE:
      return iree_hal_replay_executor_create_semaphore(executor, record);
    case IREE_HAL_REPLAY_OPERATION_CODE_ALLOCATOR_ALLOCATE_BUFFER:
      return iree_hal_replay_executor_allocate_buffer(executor, record);
    case IREE_HAL_REPLAY_OPERATION_CODE_ALLOCATOR_IMPORT_BUFFER:
      return iree_hal_replay_executor_import_buffer(executor, record);
    case IREE_HAL_REPLAY_OPERATION_CODE_DEVICE_QUEUE_ALLOCA:
      return iree_hal_replay_executor_queue_alloca(executor, record);
    case IREE_HAL_REPLAY_OPERATION_CODE_DEVICE_QUEUE_DEALLOCA:
      return iree_hal_replay_executor_queue_dealloca(executor, record);
    case IREE_HAL_REPLAY_OPERATION_CODE_DEVICE_QUEUE_FILL:
      return iree_hal_replay_executor_queue_fill(executor, record);
    case IREE_HAL_REPLAY_OPERATION_CODE_DEVICE_QUEUE_UPDATE:
      return iree_hal_replay_executor_queue_update(executor, record);
    case IREE_HAL_REPLAY_OPERATION_CODE_DEVICE_QUEUE_COPY:
      return iree_hal_replay_executor_queue_copy(executor, record);
    case IREE_HAL_REPLAY_OPERATION_CODE_DEVICE_QUEUE_READ:
      return iree_hal_replay_executor_queue_read(executor, record);
    case IREE_HAL_REPLAY_OPERATION_CODE_DEVICE_QUEUE_WRITE:
      return iree_hal_replay_executor_queue_write(executor, record);
    case IREE_HAL_REPLAY_OPERATION_CODE_BUFFER_FLUSH_RANGE:
    case IREE_HAL_REPLAY_OPERATION_CODE_BUFFER_UNMAP_RANGE:
      if (record->header.payload_type ==
          IREE_HAL_REPLAY_PAYLOAD_TYPE_BUFFER_RANGE_DATA) {
        return iree_hal_replay_executor_replay_buffer_range_data(executor,
                                                                 record);
      }
      return iree_ok_status();
    case IREE_HAL_REPLAY_OPERATION_CODE_COMMAND_BUFFER_BEGIN: {
      iree_hal_replay_object_entry_t* entry = NULL;
      IREE_RETURN_IF_ERROR(iree_hal_replay_executor_lookup(
          executor, record->header.object_id,
          IREE_HAL_REPLAY_OBJECT_TYPE_COMMAND_BUFFER, &entry));
      return iree_hal_command_buffer_begin(entry->value.command_buffer);
    }
    case IREE_HAL_REPLAY_OPERATION_CODE_COMMAND_BUFFER_END: {
      iree_hal_replay_object_entry_t* entry = NULL;
      IREE_RETURN_IF_ERROR(iree_hal_replay_executor_lookup(
          executor, record->header.object_id,
          IREE_HAL_REPLAY_OBJECT_TYPE_COMMAND_BUFFER, &entry));
      return iree_hal_command_buffer_end(entry->value.command_buffer);
    }
    case IREE_HAL_REPLAY_OPERATION_CODE_COMMAND_BUFFER_EXECUTION_BARRIER:
      return iree_hal_replay_executor_command_buffer_execution_barrier(executor,
                                                                       record);
    case IREE_HAL_REPLAY_OPERATION_CODE_COMMAND_BUFFER_SIGNAL_EVENT:
    case IREE_HAL_REPLAY_OPERATION_CODE_COMMAND_BUFFER_RESET_EVENT:
      return iree_hal_replay_executor_command_buffer_event(executor, record);
    case IREE_HAL_REPLAY_OPERATION_CODE_COMMAND_BUFFER_WAIT_EVENTS:
      return iree_hal_replay_executor_command_buffer_wait_events(executor,
                                                                 record);
    case IREE_HAL_REPLAY_OPERATION_CODE_COMMAND_BUFFER_DISPATCH:
      return iree_hal_replay_executor_command_buffer_dispatch(executor, record);
    case IREE_HAL_REPLAY_OPERATION_CODE_COMMAND_BUFFER_FILL_BUFFER:
      return iree_hal_replay_executor_command_buffer_fill_buffer(executor,
                                                                 record);
    case IREE_HAL_REPLAY_OPERATION_CODE_COMMAND_BUFFER_UPDATE_BUFFER:
      return iree_hal_replay_executor_command_buffer_update_buffer(executor,
                                                                   record);
    case IREE_HAL_REPLAY_OPERATION_CODE_COMMAND_BUFFER_COPY_BUFFER:
      return iree_hal_replay_executor_command_buffer_copy_buffer(executor,
                                                                 record);
    case IREE_HAL_REPLAY_OPERATION_CODE_DEVICE_QUEUE_DISPATCH:
      return iree_hal_replay_executor_queue_dispatch(executor, record);
    case IREE_HAL_REPLAY_OPERATION_CODE_DEVICE_QUEUE_EXECUTE:
      return iree_hal_replay_executor_queue_execute(executor, record);
    case IREE_HAL_REPLAY_OPERATION_CODE_DEVICE_QUEUE_FLUSH: {
      iree_hal_replay_object_entry_t* entry = NULL;
      IREE_RETURN_IF_ERROR(iree_hal_replay_executor_lookup(
          executor, record->header.object_id,
          IREE_HAL_REPLAY_OBJECT_TYPE_DEVICE, &entry));
      return iree_hal_device_queue_flush(entry->value.device,
                                         IREE_HAL_QUEUE_AFFINITY_ANY);
    }
    default:
      return iree_make_status(
          IREE_STATUS_UNIMPLEMENTED, "replay operation %s is not implemented",
          iree_hal_replay_operation_code_string(record->header.operation_code));
  }
}

static iree_status_t iree_hal_replay_executor_replay_unsupported(
    const iree_hal_replay_file_record_t* record) {
  if (record->header.status_code != IREE_STATUS_OK) {
    return iree_ok_status();
  }
  return iree_make_status(
      IREE_STATUS_UNIMPLEMENTED,
      "replay contains unsupported captured operation %s",
      iree_hal_replay_operation_code_string(record->header.operation_code));
}

static iree_status_t iree_hal_replay_plan_execute_record(
    iree_hal_replay_executor_t* executor,
    const iree_hal_replay_plan_record_t* plan_record) {
  const iree_hal_replay_file_record_t* record = &plan_record->file_record;
  switch (record->header.record_type) {
    case IREE_HAL_REPLAY_FILE_RECORD_TYPE_SESSION:
      return iree_ok_status();
    case IREE_HAL_REPLAY_FILE_RECORD_TYPE_OBJECT:
      return iree_hal_replay_executor_replay_object(executor, record);
    case IREE_HAL_REPLAY_FILE_RECORD_TYPE_OPERATION:
      switch (plan_record->kind) {
        case IREE_HAL_REPLAY_PLAN_RECORD_KIND_SKIP:
          return iree_ok_status();
        case IREE_HAL_REPLAY_PLAN_RECORD_KIND_SCOPE:
          return iree_hal_replay_plan_execute_scope(executor, plan_record);
        case IREE_HAL_REPLAY_PLAN_RECORD_KIND_COMMAND_BUFFER_EXECUTION_BARRIER:
          return iree_hal_replay_plan_execute_command_buffer_execution_barrier(
              executor, &plan_record->payload.command_buffer_execution_barrier);
        case IREE_HAL_REPLAY_PLAN_RECORD_KIND_COMMAND_BUFFER_DISPATCH:
          return iree_hal_replay_plan_execute_command_buffer_dispatch(
              executor, &plan_record->payload.command_buffer_dispatch);
        case IREE_HAL_REPLAY_PLAN_RECORD_KIND_QUEUE_EXECUTE:
          return iree_hal_replay_plan_execute_queue_execute(
              executor, &plan_record->payload.queue_execute);
        case IREE_HAL_REPLAY_PLAN_RECORD_KIND_GENERIC:
        default:
          return iree_hal_replay_executor_replay_operation(executor, record);
      }
    case IREE_HAL_REPLAY_FILE_RECORD_TYPE_UNSUPPORTED:
      return iree_hal_replay_executor_replay_unsupported(record);
    default:
      return iree_make_status(
          IREE_STATUS_UNIMPLEMENTED, "replay record type %s is not executable",
          iree_hal_replay_file_record_type_string(record->header.record_type));
  }
}

IREE_API_EXPORT iree_status_t iree_hal_replay_plan_execute(
    const iree_hal_replay_plan_t* plan, iree_hal_device_group_t* device_group,
    const iree_hal_replay_execute_options_t* options,
    iree_allocator_t host_allocator) {
  IREE_ASSERT_ARGUMENT(plan);
  IREE_ASSERT_ARGUMENT(device_group);
  iree_hal_replay_execute_options_t default_options =
      iree_hal_replay_execute_options_default();
  if (!options) options = &default_options;
  if (IREE_UNLIKELY(options->flags != IREE_HAL_REPLAY_EXECUTE_FLAG_NONE)) {
    return iree_make_status(IREE_STATUS_INVALID_ARGUMENT,
                            "replay execute flags are unsupported");
  }
  if (IREE_UNLIKELY(options->file_path_remap_count != 0 &&
                    !options->file_path_remaps)) {
    return iree_make_status(
        IREE_STATUS_INVALID_ARGUMENT,
        "replay execute file path remaps require a remap list");
  }

  iree_hal_replay_executor_t executor;
  IREE_RETURN_IF_ERROR(iree_hal_replay_executor_initialize(
      &executor, plan, device_group, options, host_allocator));

  iree_status_t status = iree_ok_status();
  for (iree_host_size_t i = 0;
       i < plan->record_count && iree_status_is_ok(status); ++i) {
    status = iree_hal_replay_plan_execute_record(&executor, &plan->records[i]);
  }
  if (iree_status_is_ok(status) &&
      executor.next_device_index !=
          iree_hal_device_group_device_count(executor.device_group)) {
    status = iree_make_status(
        IREE_STATUS_FAILED_PRECONDITION,
        "replay captured %" PRIhsz
        " device(s) but the target device group contains %" PRIhsz " device(s)",
        executor.next_device_index,
        iree_hal_device_group_device_count(executor.device_group));
  }

  iree_hal_replay_executor_deinitialize(&executor);
  return status;
}

IREE_API_EXPORT iree_status_t iree_hal_replay_execute_file(
    iree_const_byte_span_t file_contents, iree_hal_device_group_t* device_group,
    const iree_hal_replay_execute_options_t* options,
    iree_allocator_t host_allocator) {
  iree_hal_replay_plan_t* plan = NULL;
  IREE_RETURN_IF_ERROR(
      iree_hal_replay_plan_create(file_contents, host_allocator, &plan));
  iree_status_t status =
      iree_hal_replay_plan_execute(plan, device_group, options, host_allocator);
  iree_hal_replay_plan_destroy(plan);
  return status;
}
