// 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/tooling/profile/projection.h"

#include <string.h>

#include "iree/tooling/profile/reader.h"

static bool iree_profile_command_operation_filter_matches(
    iree_string_view_t operation_name, iree_string_view_t key,
    iree_string_view_t filter) {
  return iree_profile_key_matches(operation_name, filter) ||
         iree_profile_key_matches(key, filter);
}

static bool iree_profile_projection_try_fit_driver_host_cpu_clock(
    const iree_profile_dispatch_context_t* context,
    uint32_t physical_device_ordinal,
    iree_profile_model_clock_fit_t* out_clock_fit) {
  const iree_profile_model_device_t* device =
      iree_profile_model_find_device(&context->model, physical_device_ordinal);
  return iree_profile_model_device_try_fit_clock_exact(
      device, IREE_PROFILE_MODEL_CLOCK_TIME_DOMAIN_HOST_CPU_TIMESTAMP_NS,
      out_clock_fit);
}

typedef struct iree_profile_projection_dispatch_timing_t {
  // True when the aggregate's device has a host CPU clock fit.
  bool has_clock_fit;
  // True when min/max/total tick values were converted to nanoseconds.
  bool has_time_ns;
  // Average dispatch duration in device ticks.
  double average_ticks;
  // Nanoseconds per device tick when |has_clock_fit| is true.
  double ns_per_tick;
  // Minimum dispatch duration in nanoseconds when |has_time_ns| is true.
  int64_t minimum_ns;
  // Maximum dispatch duration in nanoseconds when |has_time_ns| is true.
  int64_t maximum_ns;
  // Total dispatch duration in nanoseconds when |has_time_ns| is true.
  int64_t total_ns;
} iree_profile_projection_dispatch_timing_t;

static iree_profile_projection_dispatch_timing_t
iree_profile_projection_calculate_dispatch_timing(
    const iree_profile_dispatch_context_t* context,
    const iree_profile_dispatch_aggregate_t* aggregate) {
  iree_profile_projection_dispatch_timing_t timing = {0};
  timing.average_ticks =
      aggregate->valid_count
          ? aggregate->total_ticks / (double)aggregate->valid_count
          : 0.0;

  iree_profile_model_clock_fit_t clock_fit;
  timing.has_clock_fit = iree_profile_projection_try_fit_driver_host_cpu_clock(
      context, aggregate->physical_device_ordinal, &clock_fit);
  timing.ns_per_tick =
      timing.has_clock_fit
          ? iree_profile_model_clock_fit_ns_per_tick(&clock_fit)
          : 0.0;
  timing.has_time_ns =
      timing.has_clock_fit && aggregate->valid_count != 0 &&
      iree_profile_model_clock_fit_scale_ticks_to_ns(
          &clock_fit, aggregate->minimum_ticks, &timing.minimum_ns) &&
      iree_profile_model_clock_fit_scale_ticks_to_ns(
          &clock_fit, aggregate->maximum_ticks, &timing.maximum_ns) &&
      iree_profile_model_clock_fit_scale_ticks_to_ns(
          &clock_fit, aggregate->total_ticks, &timing.total_ns);
  return timing;
}

typedef struct iree_profile_projection_submission_timing_t {
  // True when the aggregate's device has a host CPU clock fit.
  bool has_clock_fit;
  // True when |total_dispatch_ticks| was converted to nanoseconds.
  bool has_total_dispatch_ns;
  // Span from earliest dispatch start to latest dispatch end in device ticks.
  double span_ticks;
  // Nanoseconds per device tick when |has_clock_fit| is true.
  double ns_per_tick;
  // Total dispatch duration in nanoseconds when |has_total_dispatch_ns| is
  // true.
  int64_t total_dispatch_ns;
} iree_profile_projection_submission_timing_t;

static iree_profile_projection_submission_timing_t
iree_profile_projection_calculate_submission_timing(
    const iree_profile_dispatch_context_t* context,
    uint32_t physical_device_ordinal, uint64_t earliest_start_tick,
    uint64_t latest_end_tick, uint64_t valid_count,
    uint64_t total_dispatch_ticks) {
  iree_profile_projection_submission_timing_t timing = {0};
  timing.span_ticks =
      iree_profile_model_span_ticks(earliest_start_tick, latest_end_tick);

  iree_profile_model_clock_fit_t clock_fit;
  timing.has_clock_fit = iree_profile_projection_try_fit_driver_host_cpu_clock(
      context, physical_device_ordinal, &clock_fit);
  timing.ns_per_tick =
      timing.has_clock_fit
          ? iree_profile_model_clock_fit_ns_per_tick(&clock_fit)
          : 0.0;
  timing.has_total_dispatch_ns =
      timing.has_clock_fit && valid_count != 0 &&
      iree_profile_model_clock_fit_scale_ticks_to_ns(
          &clock_fit, total_dispatch_ticks, &timing.total_dispatch_ns);
  return timing;
}

static double iree_profile_projection_dispatch_stddev_ticks(
    const iree_profile_dispatch_aggregate_t* aggregate) {
  const double variance_ticks =
      aggregate->valid_count > 1
          ? aggregate->m2_ticks / (double)(aggregate->valid_count - 1)
          : 0.0;
  return iree_profile_sqrt_f64(variance_ticks);
}

static double iree_profile_projection_host_dispatch_stddev_ns(
    const iree_profile_host_dispatch_aggregate_t* aggregate) {
  const double variance_ns =
      aggregate->valid_count > 1
          ? aggregate->m2_ns / (double)(aggregate->valid_count - 1)
          : 0.0;
  return iree_profile_sqrt_f64(variance_ns);
}

static void iree_profile_dispatch_print_event_jsonl(
    const iree_profile_dispatch_event_row_t* row, FILE* file) {
  const iree_hal_profile_file_record_t* file_record = row->file_record;
  const iree_hal_profile_dispatch_event_t* event = row->event;
  const bool is_valid = event->start_tick != 0 && event->end_tick != 0 &&
                        event->end_tick >= event->start_tick;
  const uint64_t duration_ticks =
      is_valid ? event->end_tick - event->start_tick : 0;
  int64_t duration_ns = 0;
  const bool has_duration_ns =
      row->has_clock_fit && is_valid &&
      iree_profile_model_clock_fit_scale_ticks_to_ns(
          row->clock_fit, duration_ticks, &duration_ns);

  fprintf(file,
          "{\"type\":\"dispatch_event\",\"physical_device_ordinal\":%u"
          ",\"queue_ordinal\":%u,\"stream_id\":%" PRIu64
          ",\"event_id\":%" PRIu64 ",\"submission_id\":%" PRIu64,
          file_record->header.physical_device_ordinal,
          file_record->header.queue_ordinal, file_record->header.stream_id,
          event->event_id, event->submission_id);
  fprintf(file,
          ",\"command_buffer_id\":%" PRIu64
          ",\"command_index\":%u"
          ",\"executable_id\":%" PRIu64 ",\"function_ordinal\":%u",
          event->command_buffer_id, event->command_index, event->executable_id,
          event->function_ordinal);
  fprintf(file, ",\"key\":");
  iree_profile_fprint_json_string(file, row->key);
  fprintf(file,
          ",\"flags\":%u,\"workgroup_count\":[%u,%u,%u]"
          ",\"workgroup_size\":[%u,%u,%u]",
          event->flags, event->workgroup_count[0], event->workgroup_count[1],
          event->workgroup_count[2], event->workgroup_size[0],
          event->workgroup_size[1], event->workgroup_size[2]);
  fprintf(file,
          ",\"start_tick\":%" PRIu64 ",\"end_tick\":%" PRIu64
          ",\"duration_ticks\":%" PRIu64 ",\"valid\":%s",
          event->start_tick, event->end_tick, duration_ticks,
          is_valid ? "true" : "false");
  fprintf(file, ",\"clock_fit_available\":%s",
          row->has_clock_fit ? "true" : "false");
  fprintf(file,
          ",\"device_tick_domain\":\"device_tick\""
          ",\"duration_time_domain\":\"device_tick_duration_ns\""
          ",\"duration_ns\":%" PRId64,
          has_duration_ns ? duration_ns : 0);
  fputs("}\n", file);
}

static void iree_profile_dispatch_print_host_event_jsonl(
    const iree_profile_host_dispatch_event_row_t* row, FILE* file) {
  const iree_hal_profile_host_execution_event_t* event = row->event;
  const bool is_valid = event->start_host_time_ns >= 0 &&
                        event->end_host_time_ns >= event->start_host_time_ns;
  const int64_t duration_ns =
      is_valid ? event->end_host_time_ns - event->start_host_time_ns : 0;

  fprintf(file,
          "{\"type\":\"host_dispatch_event\",\"timing_source\":"
          "\"host_execution_event\",\"time_domain\":\"iree_host_time_ns\""
          ",\"duration_time_domain\":\"iree_host_duration_ns\""
          ",\"physical_device_ordinal\":%u"
          ",\"queue_ordinal\":%u,\"stream_id\":%" PRIu64
          ",\"event_id\":%" PRIu64 ",\"submission_id\":%" PRIu64,
          event->physical_device_ordinal, event->queue_ordinal,
          event->stream_id, event->event_id, event->submission_id);
  fprintf(file,
          ",\"command_buffer_id\":%" PRIu64
          ",\"command_index\":%u"
          ",\"executable_id\":%" PRIu64 ",\"function_ordinal\":%u",
          event->command_buffer_id, event->command_index, event->executable_id,
          event->function_ordinal);
  fprintf(file, ",\"key\":");
  iree_profile_fprint_json_string(file, row->key);
  fprintf(file,
          ",\"flags\":%u,\"status_code\":%u"
          ",\"workgroup_count\":[%u,%u,%u]"
          ",\"workgroup_size\":[%u,%u,%u]",
          event->flags, event->status_code, event->workgroup_count[0],
          event->workgroup_count[1], event->workgroup_count[2],
          event->workgroup_size[0], event->workgroup_size[1],
          event->workgroup_size[2]);
  fprintf(file,
          ",\"start_host_time_ns\":%" PRId64 ",\"end_host_time_ns\":%" PRId64
          ",\"duration_ns\":%" PRId64
          ",\"valid\":%s"
          ",\"tile_count\":%" PRIu64 ",\"tile_duration_sum_ns\":%" PRId64 "}\n",
          event->start_host_time_ns, event->end_host_time_ns, duration_ns,
          is_valid ? "true" : "false", event->tile_count,
          event->tile_duration_sum_ns);
}

static iree_status_t iree_profile_projection_emit_dispatch_event(
    void* user_data, const iree_profile_dispatch_event_row_t* row) {
  FILE* file = (FILE*)user_data;
  iree_profile_dispatch_print_event_jsonl(row, file);
  return iree_ok_status();
}

static iree_status_t iree_profile_projection_emit_host_dispatch_event(
    void* user_data, const iree_profile_host_dispatch_event_row_t* row) {
  FILE* file = (FILE*)user_data;
  iree_profile_dispatch_print_host_event_jsonl(row, file);
  return iree_ok_status();
}

static iree_status_t iree_profile_command_count_matching_operations(
    const iree_profile_dispatch_context_t* context, iree_string_view_t filter,
    int64_t id_filter, iree_host_size_t* out_operation_count) {
  *out_operation_count = 0;
  iree_host_size_t operation_count = 0;
  iree_status_t status = iree_ok_status();

  if (id_filter >= 0) {
    const iree_profile_model_command_buffer_t* command_buffer =
        iree_profile_model_find_command_buffer(&context->model,
                                               (uint64_t)id_filter);
    iree_host_size_t operation_index =
        command_buffer ? command_buffer->first_operation_index
                       : IREE_HOST_SIZE_MAX;
    while (operation_index != IREE_HOST_SIZE_MAX && iree_status_is_ok(status)) {
      const iree_hal_profile_command_operation_record_t* operation =
          &context->model.command_operations[operation_index].record;
      const char* operation_name =
          iree_profile_command_operation_type_name(operation->type);
      char numeric_buffer[128];
      iree_string_view_t key = iree_string_view_empty();
      status = iree_profile_model_resolve_command_operation_key(
          &context->model, operation, numeric_buffer, sizeof(numeric_buffer),
          &key);
      if (iree_status_is_ok(status) &&
          iree_profile_command_operation_filter_matches(
              iree_make_cstring_view(operation_name), key, filter)) {
        ++operation_count;
      }
      operation_index = context->model.command_operations[operation_index]
                            .next_operation_index;
    }
    if (iree_status_is_ok(status)) {
      *out_operation_count = operation_count;
    }
    return status;
  }

  for (iree_host_size_t i = 0;
       i < context->model.command_operation_count && iree_status_is_ok(status);
       ++i) {
    const iree_hal_profile_command_operation_record_t* operation =
        &context->model.command_operations[i].record;
    const char* operation_name =
        iree_profile_command_operation_type_name(operation->type);
    char numeric_buffer[128];
    iree_string_view_t key = iree_string_view_empty();
    status = iree_profile_model_resolve_command_operation_key(
        &context->model, operation, numeric_buffer, sizeof(numeric_buffer),
        &key);
    if (iree_status_is_ok(status) &&
        iree_profile_command_operation_filter_matches(
            iree_make_cstring_view(operation_name), key, filter)) {
      ++operation_count;
    }
  }
  if (iree_status_is_ok(status)) {
    *out_operation_count = operation_count;
  }
  return status;
}

static iree_status_t iree_profile_command_print_operation_text(
    const iree_profile_dispatch_context_t* context,
    const iree_hal_profile_command_operation_record_t* operation,
    iree_string_view_t filter, FILE* file) {
  const char* operation_name =
      iree_profile_command_operation_type_name(operation->type);
  char numeric_buffer[128];
  iree_string_view_t key = iree_string_view_empty();
  IREE_RETURN_IF_ERROR(iree_profile_model_resolve_command_operation_key(
      &context->model, operation, numeric_buffer, sizeof(numeric_buffer),
      &key));
  if (!iree_profile_command_operation_filter_matches(
          iree_make_cstring_view(operation_name), key, filter)) {
    return iree_ok_status();
  }

  fprintf(file, "    command[%u]: op=%s flags=0x%x", operation->command_index,
          operation_name, operation->flags);
  if (iree_hal_profile_command_operation_has_block_structure(operation)) {
    fprintf(file, " block=%u local=%u", operation->block_ordinal,
            operation->block_command_ordinal);
  }
  if (operation->type == IREE_HAL_PROFILE_COMMAND_OPERATION_TYPE_DISPATCH) {
    fprintf(file,
            " executable=%" PRIu64
            " function=%u key=%.*s bindings=%u"
            " workgroups=[%u,%u,%u] workgroup_size=[%u,%u,%u]",
            operation->executable_id, operation->function_ordinal,
            (int)key.size, key.data, operation->binding_count,
            operation->workgroup_count[0], operation->workgroup_count[1],
            operation->workgroup_count[2], operation->workgroup_size[0],
            operation->workgroup_size[1], operation->workgroup_size[2]);
  } else if (operation->type == IREE_HAL_PROFILE_COMMAND_OPERATION_TYPE_FILL ||
             operation->type ==
                 IREE_HAL_PROFILE_COMMAND_OPERATION_TYPE_UPDATE) {
    fprintf(file, " target=%u target_offset=%" PRIu64 " length=%" PRIu64,
            operation->target_ordinal, operation->target_offset,
            operation->length);
  } else if (operation->type == IREE_HAL_PROFILE_COMMAND_OPERATION_TYPE_COPY) {
    fprintf(file,
            " source=%u source_offset=%" PRIu64
            " target=%u target_offset=%" PRIu64 " length=%" PRIu64,
            operation->source_ordinal, operation->source_offset,
            operation->target_ordinal, operation->target_offset,
            operation->length);
  } else if (operation->type ==
                 IREE_HAL_PROFILE_COMMAND_OPERATION_TYPE_BRANCH ||
             operation->type ==
                 IREE_HAL_PROFILE_COMMAND_OPERATION_TYPE_COND_BRANCH) {
    fprintf(file, " target_block=%u alternate_block=%u",
            operation->target_block_ordinal,
            operation->alternate_block_ordinal);
  }
  fputc('\n', file);
  return iree_ok_status();
}

static iree_status_t iree_profile_command_print_operation_jsonl(
    const iree_profile_dispatch_context_t* context,
    const iree_hal_profile_command_operation_record_t* operation,
    iree_string_view_t filter, FILE* file) {
  const char* operation_name =
      iree_profile_command_operation_type_name(operation->type);
  char numeric_buffer[128];
  iree_string_view_t key = iree_string_view_empty();
  IREE_RETURN_IF_ERROR(iree_profile_model_resolve_command_operation_key(
      &context->model, operation, numeric_buffer, sizeof(numeric_buffer),
      &key));
  if (!iree_profile_command_operation_filter_matches(
          iree_make_cstring_view(operation_name), key, filter)) {
    return iree_ok_status();
  }

  fprintf(file,
          "{\"type\":\"command_operation\",\"command_buffer_id\":%" PRIu64
          ",\"command_index\":%u,\"op\":\"%s\",\"flags\":%u"
          ",\"block_structure\":%s"
          ",\"block_ordinal\":%u,\"block_command_ordinal\":%u",
          operation->command_buffer_id, operation->command_index,
          operation_name, operation->flags,
          iree_hal_profile_command_operation_has_block_structure(operation)
              ? "true"
              : "false",
          operation->block_ordinal, operation->block_command_ordinal);
  fprintf(file, ",\"key\":");
  iree_profile_fprint_json_string(file, key);
  fprintf(file,
          ",\"executable_id\":%" PRIu64
          ",\"function_ordinal\":%u"
          ",\"binding_count\":%u,\"workgroup_count\":[%u,%u,%u]"
          ",\"workgroup_size\":[%u,%u,%u]"
          ",\"source_ordinal\":%u,\"target_ordinal\":%u"
          ",\"source_offset\":%" PRIu64 ",\"target_offset\":%" PRIu64
          ",\"length\":%" PRIu64
          ",\"target_block_ordinal\":%u,\"alternate_block_ordinal\":%u}\n",
          operation->executable_id, operation->function_ordinal,
          operation->binding_count, operation->workgroup_count[0],
          operation->workgroup_count[1], operation->workgroup_count[2],
          operation->workgroup_size[0], operation->workgroup_size[1],
          operation->workgroup_size[2], operation->source_ordinal,
          operation->target_ordinal, operation->source_offset,
          operation->target_offset, operation->length,
          operation->target_block_ordinal, operation->alternate_block_ordinal);
  return iree_ok_status();
}

static void iree_profile_dispatch_print_text_header(
    const iree_profile_dispatch_context_t* context, iree_string_view_t filter,
    FILE* file) {
  fprintf(file, "IREE HAL profile dispatch summary\n");
  fprintf(file, "filter: %.*s\n", (int)filter.size, filter.data);
  fprintf(file,
          "dispatches: total=%" PRIu64 " matched=%" PRIu64 " valid=%" PRIu64
          " invalid=%" PRIu64 " groups=%" PRIhsz "\n",
          context->total_dispatch_count, context->matched_dispatch_count,
          context->valid_dispatch_count, context->invalid_dispatch_count,
          context->aggregate_count);
  fprintf(
      file,
      "host_dispatches: total=%" PRIu64 " matched=%" PRIu64 " valid=%" PRIu64
      " invalid=%" PRIu64 " groups=%" PRIhsz "\n",
      context->total_host_dispatch_count, context->matched_host_dispatch_count,
      context->valid_host_dispatch_count, context->invalid_host_dispatch_count,
      context->host_dispatch_aggregate_count);
  fprintf(file, "groups:\n");
}

static void iree_profile_dispatch_print_text_group(
    const iree_profile_dispatch_context_t* context,
    const iree_profile_dispatch_aggregate_t* aggregate, iree_string_view_t key,
    FILE* file) {
  const iree_profile_projection_dispatch_timing_t timing =
      iree_profile_projection_calculate_dispatch_timing(context, aggregate);
  const double stddev_ticks =
      iree_profile_projection_dispatch_stddev_ticks(aggregate);
  fprintf(file, "  %.*s\n", (int)key.size, key.data);
  fprintf(file,
          "    device=%u executable=%" PRIu64 " function=%u count=%" PRIu64
          " valid=%" PRIu64 " invalid=%" PRIu64 "\n",
          aggregate->physical_device_ordinal, aggregate->executable_id,
          aggregate->function_ordinal, aggregate->dispatch_count,
          aggregate->valid_count, aggregate->invalid_count);
  if (aggregate->valid_count != 0) {
    fprintf(file,
            "    ticks: min=%" PRIu64 " avg=%.3f stddev=%.3f max=%" PRIu64
            " total=%" PRIu64 "\n",
            aggregate->minimum_ticks, aggregate->mean_ticks, stddev_ticks,
            aggregate->maximum_ticks, aggregate->total_ticks);
    if (timing.has_time_ns) {
      fprintf(file,
              "    time_ns: min=%" PRId64
              " avg=%.3f stddev=%.3f"
              " max=%" PRId64 " total=%" PRId64 "\n",
              timing.minimum_ns, aggregate->mean_ticks * timing.ns_per_tick,
              stddev_ticks * timing.ns_per_tick, timing.maximum_ns,
              timing.total_ns);
    } else {
      fprintf(file, "    time_ns: unavailable\n");
    }
  }
  fprintf(file,
          "    last_geometry: workgroup_count=%ux%ux%u"
          " workgroup_size=%ux%ux%u\n",
          aggregate->last_workgroup_count[0],
          aggregate->last_workgroup_count[1],
          aggregate->last_workgroup_count[2], aggregate->last_workgroup_size[0],
          aggregate->last_workgroup_size[1], aggregate->last_workgroup_size[2]);
}

static void iree_profile_dispatch_print_text_host_group(
    const iree_profile_host_dispatch_aggregate_t* aggregate,
    iree_string_view_t key, FILE* file) {
  const double stddev_ns =
      iree_profile_projection_host_dispatch_stddev_ns(aggregate);
  fprintf(file, "  %.*s\n", (int)key.size, key.data);
  fprintf(file,
          "    timing=host_execution device=%u executable=%" PRIu64
          " function=%u count=%" PRIu64 " valid=%" PRIu64 " invalid=%" PRIu64
          "\n",
          aggregate->physical_device_ordinal, aggregate->executable_id,
          aggregate->function_ordinal, aggregate->dispatch_count,
          aggregate->valid_count, aggregate->invalid_count);
  if (aggregate->valid_count != 0) {
    fprintf(file,
            "    time_ns: min=%" PRId64
            " avg=%.3f stddev=%.3f"
            " max=%" PRId64 " total=%" PRId64 " tile_count=%" PRIu64
            " tile_sum_ns=%" PRId64 "\n",
            aggregate->minimum_ns, aggregate->mean_ns, stddev_ns,
            aggregate->maximum_ns, aggregate->total_ns,
            aggregate->total_tile_count, aggregate->total_tile_duration_sum_ns);
  }
  fprintf(file,
          "    last_geometry: workgroup_count=%ux%ux%u"
          " workgroup_size=%ux%ux%u\n",
          aggregate->last_workgroup_count[0],
          aggregate->last_workgroup_count[1],
          aggregate->last_workgroup_count[2], aggregate->last_workgroup_size[0],
          aggregate->last_workgroup_size[1], aggregate->last_workgroup_size[2]);
}

static iree_status_t iree_profile_dispatch_print_text(
    const iree_profile_dispatch_context_t* context, iree_string_view_t filter,
    FILE* file) {
  iree_profile_dispatch_print_text_header(context, filter, file);

  iree_status_t status = iree_ok_status();
  for (iree_host_size_t i = 0;
       i < context->aggregate_count && iree_status_is_ok(status); ++i) {
    const iree_profile_dispatch_aggregate_t* aggregate =
        &context->aggregates[i];
    char numeric_buffer[128];
    iree_string_view_t key = iree_string_view_empty();
    status = iree_profile_model_resolve_dispatch_key(
        &context->model, aggregate->physical_device_ordinal,
        aggregate->executable_id, aggregate->function_ordinal, numeric_buffer,
        sizeof(numeric_buffer), &key);
    if (iree_status_is_ok(status)) {
      iree_profile_dispatch_print_text_group(context, aggregate, key, file);
    }
  }
  for (iree_host_size_t i = 0;
       i < context->host_dispatch_aggregate_count && iree_status_is_ok(status);
       ++i) {
    const iree_profile_host_dispatch_aggregate_t* aggregate =
        &context->host_dispatch_aggregates[i];
    char numeric_buffer[128];
    iree_string_view_t key = iree_string_view_empty();
    status = iree_profile_model_resolve_dispatch_key(
        &context->model, aggregate->physical_device_ordinal,
        aggregate->executable_id, aggregate->function_ordinal, numeric_buffer,
        sizeof(numeric_buffer), &key);
    if (iree_status_is_ok(status)) {
      iree_profile_dispatch_print_text_host_group(aggregate, key, file);
    }
  }
  return status;
}

static void iree_profile_dispatch_print_jsonl_summary(
    const iree_profile_dispatch_context_t* context, iree_string_view_t filter,
    bool emit_events, FILE* file) {
  fprintf(file, "{\"type\":\"dispatch_summary\",\"mode\":\"%s\",\"filter\":",
          emit_events ? "events" : "aggregate");
  iree_profile_fprint_json_string(file, filter);
  fprintf(file,
          ",\"total_dispatches\":%" PRIu64 ",\"matched_dispatches\":%" PRIu64
          ",\"valid_dispatches\":%" PRIu64 ",\"invalid_dispatches\":%" PRIu64
          ",\"aggregate_groups\":%" PRIhsz ",\"total_host_dispatches\":%" PRIu64
          ",\"matched_host_dispatches\":%" PRIu64
          ",\"valid_host_dispatches\":%" PRIu64
          ",\"invalid_host_dispatches\":%" PRIu64
          ",\"host_aggregate_groups\":%" PRIhsz "}\n",
          context->total_dispatch_count, context->matched_dispatch_count,
          context->valid_dispatch_count, context->invalid_dispatch_count,
          context->aggregate_count, context->total_host_dispatch_count,
          context->matched_host_dispatch_count,
          context->valid_host_dispatch_count,
          context->invalid_host_dispatch_count,
          context->host_dispatch_aggregate_count);
}

static void iree_profile_dispatch_print_jsonl_group(
    const iree_profile_dispatch_context_t* context,
    const iree_profile_dispatch_aggregate_t* aggregate, iree_string_view_t key,
    FILE* file) {
  const iree_profile_projection_dispatch_timing_t timing =
      iree_profile_projection_calculate_dispatch_timing(context, aggregate);
  const double stddev_ticks =
      iree_profile_projection_dispatch_stddev_ticks(aggregate);
  fprintf(file,
          "{\"type\":\"dispatch_group\",\"physical_device_ordinal\":%u"
          ",\"executable_id\":%" PRIu64 ",\"function_ordinal\":%u,\"key\":",
          aggregate->physical_device_ordinal, aggregate->executable_id,
          aggregate->function_ordinal);
  iree_profile_fprint_json_string(file, key);
  fprintf(file,
          ",\"count\":%" PRIu64 ",\"valid\":%" PRIu64 ",\"invalid\":%" PRIu64
          ",\"min_ticks\":%" PRIu64
          ",\"avg_ticks\":%.3f,\"stddev_ticks\":%.3f"
          ",\"max_ticks\":%" PRIu64 ",\"total_ticks\":%" PRIu64,
          aggregate->dispatch_count, aggregate->valid_count,
          aggregate->invalid_count,
          aggregate->valid_count ? aggregate->minimum_ticks : 0,
          aggregate->valid_count ? aggregate->mean_ticks : 0.0, stddev_ticks,
          aggregate->valid_count ? aggregate->maximum_ticks : 0,
          aggregate->total_ticks);
  fprintf(file, ",\"clock_fit_available\":%s",
          timing.has_clock_fit ? "true" : "false");
  fprintf(file,
          ",\"time_ns_available\":%s"
          ",\"min_ns\":%" PRId64
          ",\"avg_ns\":%.3f"
          ",\"stddev_ns\":%.3f,\"max_ns\":%" PRId64 ",\"total_ns\":%" PRId64,
          timing.has_time_ns ? "true" : "false",
          timing.has_time_ns ? timing.minimum_ns : 0,
          timing.has_time_ns ? aggregate->mean_ticks * timing.ns_per_tick : 0.0,
          timing.has_time_ns ? stddev_ticks * timing.ns_per_tick : 0.0,
          timing.has_time_ns ? timing.maximum_ns : 0,
          timing.has_time_ns ? timing.total_ns : 0);
  fprintf(file,
          ",\"last_workgroup_count\":[%u,%u,%u]"
          ",\"last_workgroup_size\":[%u,%u,%u]}\n",
          aggregate->last_workgroup_count[0],
          aggregate->last_workgroup_count[1],
          aggregate->last_workgroup_count[2], aggregate->last_workgroup_size[0],
          aggregate->last_workgroup_size[1], aggregate->last_workgroup_size[2]);
}

static void iree_profile_dispatch_print_jsonl_host_group(
    const iree_profile_host_dispatch_aggregate_t* aggregate,
    iree_string_view_t key, FILE* file) {
  const double stddev_ns =
      iree_profile_projection_host_dispatch_stddev_ns(aggregate);
  fprintf(file,
          "{\"type\":\"host_dispatch_group\",\"timing_source\":"
          "\"host_execution_event\",\"time_domain\":\"iree_host_time_ns\""
          ",\"duration_time_domain\":\"iree_host_duration_ns\""
          ",\"physical_device_ordinal\":%u"
          ",\"executable_id\":%" PRIu64 ",\"function_ordinal\":%u,\"key\":",
          aggregate->physical_device_ordinal, aggregate->executable_id,
          aggregate->function_ordinal);
  iree_profile_fprint_json_string(file, key);
  fprintf(
      file,
      ",\"count\":%" PRIu64 ",\"valid\":%" PRIu64 ",\"invalid\":%" PRIu64
      ",\"min_ns\":%" PRId64
      ",\"avg_ns\":%.3f"
      ",\"stddev_ns\":%.3f,\"max_ns\":%" PRId64 ",\"total_ns\":%" PRId64
      ",\"total_tile_count\":%" PRIu64
      ",\"total_tile_duration_sum_ns\":%" PRId64
      ",\"last_workgroup_count\":[%u,%u,%u]"
      ",\"last_workgroup_size\":[%u,%u,%u]}\n",
      aggregate->dispatch_count, aggregate->valid_count,
      aggregate->invalid_count,
      aggregate->valid_count ? aggregate->minimum_ns : 0,
      aggregate->valid_count ? aggregate->mean_ns : 0.0, stddev_ns,
      aggregate->valid_count ? aggregate->maximum_ns : 0, aggregate->total_ns,
      aggregate->total_tile_count, aggregate->total_tile_duration_sum_ns,
      aggregate->last_workgroup_count[0], aggregate->last_workgroup_count[1],
      aggregate->last_workgroup_count[2], aggregate->last_workgroup_size[0],
      aggregate->last_workgroup_size[1], aggregate->last_workgroup_size[2]);
}

static iree_status_t iree_profile_dispatch_print_jsonl_aggregates(
    const iree_profile_dispatch_context_t* context, FILE* file) {
  iree_status_t status = iree_ok_status();
  for (iree_host_size_t i = 0;
       i < context->aggregate_count && iree_status_is_ok(status); ++i) {
    const iree_profile_dispatch_aggregate_t* aggregate =
        &context->aggregates[i];
    char numeric_buffer[128];
    iree_string_view_t key = iree_string_view_empty();
    status = iree_profile_model_resolve_dispatch_key(
        &context->model, aggregate->physical_device_ordinal,
        aggregate->executable_id, aggregate->function_ordinal, numeric_buffer,
        sizeof(numeric_buffer), &key);
    if (iree_status_is_ok(status)) {
      iree_profile_dispatch_print_jsonl_group(context, aggregate, key, file);
    }
  }
  for (iree_host_size_t i = 0;
       i < context->host_dispatch_aggregate_count && iree_status_is_ok(status);
       ++i) {
    const iree_profile_host_dispatch_aggregate_t* aggregate =
        &context->host_dispatch_aggregates[i];
    char numeric_buffer[128];
    iree_string_view_t key = iree_string_view_empty();
    status = iree_profile_model_resolve_dispatch_key(
        &context->model, aggregate->physical_device_ordinal,
        aggregate->executable_id, aggregate->function_ordinal, numeric_buffer,
        sizeof(numeric_buffer), &key);
    if (iree_status_is_ok(status)) {
      iree_profile_dispatch_print_jsonl_host_group(aggregate, key, file);
    }
  }
  return status;
}

static bool iree_profile_executable_matches_id(
    const iree_hal_profile_executable_record_t* executable, int64_t id_filter) {
  return id_filter < 0 || executable->executable_id == (uint64_t)id_filter;
}

typedef struct iree_profile_projection_row_list_t {
  // First row in the owner list, or IREE_HOST_SIZE_MAX.
  iree_host_size_t first_index;
  // Last row in the owner list, or IREE_HOST_SIZE_MAX.
  iree_host_size_t last_index;
} iree_profile_projection_row_list_t;

typedef struct iree_profile_projection_function_group_t {
  // Device-timestamped dispatch aggregate rows grouped by function.
  iree_profile_projection_row_list_t device_dispatches;
  // Host-timestamped dispatch aggregate rows grouped by function.
  iree_profile_projection_row_list_t host_dispatches;
} iree_profile_projection_function_group_t;

typedef struct iree_profile_projection_command_buffer_group_t {
  // Device-timestamped command-buffer execution rows grouped by command buffer.
  iree_profile_projection_row_list_t device_executions;
  // Host-timestamped command-buffer execution rows grouped by command buffer.
  iree_profile_projection_row_list_t host_executions;
} iree_profile_projection_command_buffer_group_t;

typedef struct iree_profile_projection_queue_group_t {
  // Device dispatch submission aggregate rows grouped by queue.
  iree_profile_projection_row_list_t submissions;
  // Device-timestamped queue event rows grouped by queue.
  iree_profile_projection_row_list_t device_events;
  // Host-timestamped execution event rows grouped by queue.
  iree_profile_projection_row_list_t host_execution_events;
  // Host-timestamped queue operation event rows grouped by queue.
  iree_profile_projection_row_list_t events;
} iree_profile_projection_queue_group_t;

typedef struct iree_profile_projection_index_t {
  // Host allocator used for owner-index arrays.
  iree_allocator_t host_allocator;
  // Function owner groups parallel to iree_profile_model_t::functions.
  iree_profile_projection_function_group_t* functions;
  // Next aggregate row indexes parallel to
  // iree_profile_dispatch_context_t::aggregates.
  iree_host_size_t* aggregate_next_indices;
  // Next host aggregate row indexes parallel to
  // iree_profile_dispatch_context_t::host_dispatch_aggregates.
  iree_host_size_t* host_dispatch_next_indices;
  // Command-buffer owner groups parallel to
  // iree_profile_model_t::command_buffers.
  iree_profile_projection_command_buffer_group_t* command_buffers;
  // Next command aggregate row indexes parallel to
  // iree_profile_dispatch_context_t::command_aggregates.
  iree_host_size_t* command_next_indices;
  // Next host command aggregate row indexes parallel to
  // iree_profile_dispatch_context_t::host_command_aggregates.
  iree_host_size_t* host_command_next_indices;
  // Queue owner groups parallel to iree_profile_model_t::queues.
  iree_profile_projection_queue_group_t* queues;
  // Next queue submission row indexes parallel to
  // iree_profile_dispatch_context_t::queue_aggregates.
  iree_host_size_t* queue_next_indices;
  // Next queue device event row indexes parallel to
  // iree_profile_queue_event_query_t::queue_device_events.
  iree_host_size_t* queue_device_event_next_indices;
  // Next host execution event row indexes parallel to
  // iree_profile_queue_event_query_t::host_execution_events.
  iree_host_size_t* host_execution_event_next_indices;
  // Next queue event row indexes parallel to
  // iree_profile_queue_event_query_t::queue_events.
  iree_host_size_t* queue_event_next_indices;
} iree_profile_projection_index_t;

static void iree_profile_projection_row_list_initialize(
    iree_profile_projection_row_list_t* row_list) {
  row_list->first_index = IREE_HOST_SIZE_MAX;
  row_list->last_index = IREE_HOST_SIZE_MAX;
}

static void iree_profile_projection_row_list_append(
    iree_profile_projection_row_list_t* row_list,
    iree_host_size_t* next_indices, iree_host_size_t row_index) {
  next_indices[row_index] = IREE_HOST_SIZE_MAX;
  if (row_list->first_index == IREE_HOST_SIZE_MAX) {
    row_list->first_index = row_index;
  } else {
    next_indices[row_list->last_index] = row_index;
  }
  row_list->last_index = row_index;
}

static void iree_profile_projection_index_initialize(
    iree_allocator_t host_allocator,
    iree_profile_projection_index_t* out_index) {
  memset(out_index, 0, sizeof(*out_index));
  out_index->host_allocator = host_allocator;
}

static void iree_profile_projection_index_deinitialize(
    iree_profile_projection_index_t* index) {
  iree_allocator_free(index->host_allocator, index->host_command_next_indices);
  iree_allocator_free(index->host_allocator, index->command_next_indices);
  iree_allocator_free(index->host_allocator, index->command_buffers);
  iree_allocator_free(index->host_allocator, index->queue_event_next_indices);
  iree_allocator_free(index->host_allocator,
                      index->host_execution_event_next_indices);
  iree_allocator_free(index->host_allocator,
                      index->queue_device_event_next_indices);
  iree_allocator_free(index->host_allocator, index->queue_next_indices);
  iree_allocator_free(index->host_allocator, index->queues);
  iree_allocator_free(index->host_allocator, index->host_dispatch_next_indices);
  iree_allocator_free(index->host_allocator, index->aggregate_next_indices);
  iree_allocator_free(index->host_allocator, index->functions);
  memset(index, 0, sizeof(*index));
}

static iree_status_t iree_profile_projection_index_allocate_row_links(
    iree_allocator_t host_allocator, iree_host_size_t row_count,
    iree_host_size_t** out_next_indices) {
  *out_next_indices = NULL;
  if (row_count == 0) return iree_ok_status();
  IREE_RETURN_IF_ERROR(iree_allocator_malloc_array(
      host_allocator, row_count, sizeof((*out_next_indices)[0]),
      (void**)out_next_indices));
  for (iree_host_size_t i = 0; i < row_count; ++i) {
    (*out_next_indices)[i] = IREE_HOST_SIZE_MAX;
  }
  return iree_ok_status();
}

static iree_status_t iree_profile_projection_index_find_function_index(
    const iree_profile_model_t* model, uint32_t physical_device_ordinal,
    uint64_t executable_id, uint32_t function_ordinal,
    iree_host_size_t* out_function_index) {
  *out_function_index = IREE_HOST_SIZE_MAX;
  const iree_profile_model_function_t* function_info =
      iree_profile_model_find_function(model, executable_id, function_ordinal);
  if (!function_info) {
    return iree_make_status(
        IREE_STATUS_DATA_LOSS,
        "aggregate references missing executable function metadata "
        "device=%u executable=%" PRIu64 " function=%u",
        physical_device_ordinal, executable_id, function_ordinal);
  }
  *out_function_index = (iree_host_size_t)(function_info - model->functions);
  return iree_ok_status();
}

static iree_status_t iree_profile_projection_index_find_command_buffer_index(
    const iree_profile_model_t* model, uint64_t command_buffer_id,
    uint64_t submission_id, iree_host_size_t* out_command_buffer_index) {
  *out_command_buffer_index = IREE_HOST_SIZE_MAX;
  const iree_profile_model_command_buffer_t* command_buffer =
      iree_profile_model_find_command_buffer(model, command_buffer_id);
  if (!command_buffer) {
    return iree_make_status(
        IREE_STATUS_DATA_LOSS,
        "aggregate references missing command-buffer metadata "
        "command_buffer=%" PRIu64 " submission=%" PRIu64,
        command_buffer_id, submission_id);
  }
  *out_command_buffer_index =
      (iree_host_size_t)(command_buffer - model->command_buffers);
  return iree_ok_status();
}

static iree_status_t iree_profile_projection_index_find_queue_index(
    const iree_profile_model_t* model, uint32_t physical_device_ordinal,
    uint32_t queue_ordinal, uint64_t stream_id, uint64_t submission_id,
    iree_host_size_t* out_queue_index) {
  *out_queue_index = IREE_HOST_SIZE_MAX;
  const iree_profile_model_queue_t* queue = iree_profile_model_find_queue(
      model, physical_device_ordinal, queue_ordinal, stream_id);
  if (!queue) {
    return iree_make_status(
        IREE_STATUS_DATA_LOSS,
        "event references missing queue metadata device=%u queue=%u "
        "stream=%" PRIu64 " submission=%" PRIu64,
        physical_device_ordinal, queue_ordinal, stream_id, submission_id);
  }
  *out_queue_index = (iree_host_size_t)(queue - model->queues);
  return iree_ok_status();
}

static iree_status_t iree_profile_projection_index_build_functions(
    const iree_profile_dispatch_context_t* context,
    iree_profile_projection_index_t* index) {
  if (context->model.function_count != 0) {
    IREE_RETURN_IF_ERROR(iree_allocator_malloc_array(
        index->host_allocator, context->model.function_count,
        sizeof(index->functions[0]), (void**)&index->functions));
    for (iree_host_size_t i = 0; i < context->model.function_count; ++i) {
      iree_profile_projection_row_list_initialize(
          &index->functions[i].device_dispatches);
      iree_profile_projection_row_list_initialize(
          &index->functions[i].host_dispatches);
    }
  }

  IREE_RETURN_IF_ERROR(iree_profile_projection_index_allocate_row_links(
      index->host_allocator, context->aggregate_count,
      &index->aggregate_next_indices));
  IREE_RETURN_IF_ERROR(iree_profile_projection_index_allocate_row_links(
      index->host_allocator, context->host_dispatch_aggregate_count,
      &index->host_dispatch_next_indices));

  for (iree_host_size_t i = 0; i < context->aggregate_count; ++i) {
    const iree_profile_dispatch_aggregate_t* aggregate =
        &context->aggregates[i];
    iree_host_size_t function_index = IREE_HOST_SIZE_MAX;
    IREE_RETURN_IF_ERROR(iree_profile_projection_index_find_function_index(
        &context->model, aggregate->physical_device_ordinal,
        aggregate->executable_id, aggregate->function_ordinal,
        &function_index));
    iree_profile_projection_row_list_append(
        &index->functions[function_index].device_dispatches,
        index->aggregate_next_indices, i);
  }

  for (iree_host_size_t i = 0; i < context->host_dispatch_aggregate_count;
       ++i) {
    const iree_profile_host_dispatch_aggregate_t* aggregate =
        &context->host_dispatch_aggregates[i];
    iree_host_size_t function_index = IREE_HOST_SIZE_MAX;
    IREE_RETURN_IF_ERROR(iree_profile_projection_index_find_function_index(
        &context->model, aggregate->physical_device_ordinal,
        aggregate->executable_id, aggregate->function_ordinal,
        &function_index));
    iree_profile_projection_row_list_append(
        &index->functions[function_index].host_dispatches,
        index->host_dispatch_next_indices, i);
  }

  return iree_ok_status();
}

static iree_status_t iree_profile_projection_index_build_command_buffers(
    const iree_profile_dispatch_context_t* context,
    iree_profile_projection_index_t* index) {
  if (context->model.command_buffer_count != 0) {
    IREE_RETURN_IF_ERROR(iree_allocator_malloc_array(
        index->host_allocator, context->model.command_buffer_count,
        sizeof(index->command_buffers[0]), (void**)&index->command_buffers));
    for (iree_host_size_t i = 0; i < context->model.command_buffer_count; ++i) {
      iree_profile_projection_row_list_initialize(
          &index->command_buffers[i].device_executions);
      iree_profile_projection_row_list_initialize(
          &index->command_buffers[i].host_executions);
    }
  }

  IREE_RETURN_IF_ERROR(iree_profile_projection_index_allocate_row_links(
      index->host_allocator, context->command_aggregate_count,
      &index->command_next_indices));
  IREE_RETURN_IF_ERROR(iree_profile_projection_index_allocate_row_links(
      index->host_allocator, context->host_command_aggregate_count,
      &index->host_command_next_indices));

  for (iree_host_size_t i = 0; i < context->command_aggregate_count; ++i) {
    const iree_profile_dispatch_command_aggregate_t* aggregate =
        &context->command_aggregates[i];
    iree_host_size_t command_buffer_index = IREE_HOST_SIZE_MAX;
    IREE_RETURN_IF_ERROR(
        iree_profile_projection_index_find_command_buffer_index(
            &context->model, aggregate->command_buffer_id,
            aggregate->submission_id, &command_buffer_index));
    iree_profile_projection_row_list_append(
        &index->command_buffers[command_buffer_index].device_executions,
        index->command_next_indices, i);
  }

  for (iree_host_size_t i = 0; i < context->host_command_aggregate_count; ++i) {
    const iree_profile_host_dispatch_command_aggregate_t* aggregate =
        &context->host_command_aggregates[i];
    iree_host_size_t command_buffer_index = IREE_HOST_SIZE_MAX;
    IREE_RETURN_IF_ERROR(
        iree_profile_projection_index_find_command_buffer_index(
            &context->model, aggregate->command_buffer_id,
            aggregate->submission_id, &command_buffer_index));
    iree_profile_projection_row_list_append(
        &index->command_buffers[command_buffer_index].host_executions,
        index->host_command_next_indices, i);
  }

  return iree_ok_status();
}

static iree_status_t iree_profile_projection_index_build_queues(
    const iree_profile_dispatch_context_t* context,
    iree_profile_projection_index_t* index) {
  if (context->model.queue_count != 0) {
    IREE_RETURN_IF_ERROR(iree_allocator_malloc_array(
        index->host_allocator, context->model.queue_count,
        sizeof(index->queues[0]), (void**)&index->queues));
    for (iree_host_size_t i = 0; i < context->model.queue_count; ++i) {
      iree_profile_projection_row_list_initialize(
          &index->queues[i].submissions);
      iree_profile_projection_row_list_initialize(
          &index->queues[i].device_events);
      iree_profile_projection_row_list_initialize(
          &index->queues[i].host_execution_events);
      iree_profile_projection_row_list_initialize(&index->queues[i].events);
    }
  }

  IREE_RETURN_IF_ERROR(iree_profile_projection_index_allocate_row_links(
      index->host_allocator, context->queue_aggregate_count,
      &index->queue_next_indices));
  IREE_RETURN_IF_ERROR(iree_profile_projection_index_allocate_row_links(
      index->host_allocator, context->queue_query.queue_device_event_count,
      &index->queue_device_event_next_indices));
  IREE_RETURN_IF_ERROR(iree_profile_projection_index_allocate_row_links(
      index->host_allocator, context->queue_query.host_execution_event_count,
      &index->host_execution_event_next_indices));
  IREE_RETURN_IF_ERROR(iree_profile_projection_index_allocate_row_links(
      index->host_allocator, context->queue_query.queue_event_count,
      &index->queue_event_next_indices));

  for (iree_host_size_t i = 0; i < context->queue_aggregate_count; ++i) {
    const iree_profile_dispatch_queue_aggregate_t* aggregate =
        &context->queue_aggregates[i];
    iree_host_size_t queue_index = IREE_HOST_SIZE_MAX;
    IREE_RETURN_IF_ERROR(iree_profile_projection_index_find_queue_index(
        &context->model, aggregate->physical_device_ordinal,
        aggregate->queue_ordinal, aggregate->stream_id,
        aggregate->submission_id, &queue_index));
    iree_profile_projection_row_list_append(
        &index->queues[queue_index].submissions, index->queue_next_indices, i);
  }

  for (iree_host_size_t i = 0;
       i < context->queue_query.queue_device_event_count; ++i) {
    const iree_hal_profile_queue_device_event_t* event =
        &context->queue_query.queue_device_events[i].record;
    iree_host_size_t queue_index = IREE_HOST_SIZE_MAX;
    IREE_RETURN_IF_ERROR(iree_profile_projection_index_find_queue_index(
        &context->model, event->physical_device_ordinal, event->queue_ordinal,
        event->stream_id, event->submission_id, &queue_index));
    iree_profile_projection_row_list_append(
        &index->queues[queue_index].device_events,
        index->queue_device_event_next_indices, i);
  }

  for (iree_host_size_t i = 0;
       i < context->queue_query.host_execution_event_count; ++i) {
    const iree_hal_profile_host_execution_event_t* event =
        &context->queue_query.host_execution_events[i].record;
    iree_host_size_t queue_index = IREE_HOST_SIZE_MAX;
    IREE_RETURN_IF_ERROR(iree_profile_projection_index_find_queue_index(
        &context->model, event->physical_device_ordinal, event->queue_ordinal,
        event->stream_id, event->submission_id, &queue_index));
    iree_profile_projection_row_list_append(
        &index->queues[queue_index].host_execution_events,
        index->host_execution_event_next_indices, i);
  }

  for (iree_host_size_t i = 0; i < context->queue_query.queue_event_count;
       ++i) {
    const iree_hal_profile_queue_event_t* event =
        &context->queue_query.queue_events[i].record;
    iree_host_size_t queue_index = IREE_HOST_SIZE_MAX;
    IREE_RETURN_IF_ERROR(iree_profile_projection_index_find_queue_index(
        &context->model, event->physical_device_ordinal, event->queue_ordinal,
        event->stream_id, event->submission_id, &queue_index));
    iree_profile_projection_row_list_append(&index->queues[queue_index].events,
                                            index->queue_event_next_indices, i);
  }

  return iree_ok_status();
}

static void iree_profile_executable_print_text_header(
    const iree_profile_dispatch_context_t* context, iree_string_view_t filter,
    FILE* file) {
  fprintf(file, "IREE HAL profile executable summary\n");
  fprintf(file, "filter: %.*s\n", (int)filter.size, filter.data);
  fprintf(file,
          "executables=%" PRIhsz " functions=%" PRIhsz
          " matched_dispatches=%" PRIu64 " valid=%" PRIu64 " invalid=%" PRIu64
          " matched_host_dispatches=%" PRIu64 " host_valid=%" PRIu64
          " host_invalid=%" PRIu64 "\n",
          context->model.executable_count, context->model.function_count,
          context->matched_dispatch_count, context->valid_dispatch_count,
          context->invalid_dispatch_count, context->matched_host_dispatch_count,
          context->valid_host_dispatch_count,
          context->invalid_host_dispatch_count);
}

static void iree_profile_executable_print_text_executable(
    const iree_hal_profile_executable_record_t* executable, FILE* file) {
  const bool has_code_object_hash = iree_all_bits_set(
      executable->flags, IREE_HAL_PROFILE_EXECUTABLE_FLAG_CODE_OBJECT_HASH);
  fprintf(file, "executable %" PRIu64 ": functions=%u flags=%u ",
          executable->executable_id, executable->function_count,
          executable->flags);
  fprintf(file, "code_object_hash=");
  if (has_code_object_hash) {
    iree_profile_fprint_hash_hex(file, executable->code_object_hash);
  } else {
    fprintf(file, "unavailable");
  }
  fputc('\n', file);
}

static void iree_profile_executable_print_text_function(
    const iree_profile_model_function_t* function_info, iree_string_view_t key,
    FILE* file) {
  const bool has_function_hash = iree_all_bits_set(
      function_info->flags,
      IREE_HAL_PROFILE_EXECUTABLE_FUNCTION_FLAG_FUNCTION_HASH);
  fprintf(file,
          "  function %u: %.*s flags=%u constants=%u bindings=%u "
          "parameters=%u workgroup_size=%ux%ux%u function_hash=",
          function_info->function_ordinal, (int)key.size, key.data,
          function_info->flags, function_info->constant_count,
          function_info->binding_count, function_info->parameter_count,
          function_info->workgroup_size[0], function_info->workgroup_size[1],
          function_info->workgroup_size[2]);
  if (has_function_hash) {
    iree_profile_fprint_hash_hex(file, function_info->function_hash);
  } else {
    fprintf(file, "unavailable");
  }
  fputc('\n', file);
}

static void iree_profile_executable_print_text_dispatch_group(
    const iree_profile_dispatch_context_t* context,
    const iree_profile_dispatch_aggregate_t* aggregate, FILE* file) {
  const iree_profile_projection_dispatch_timing_t timing =
      iree_profile_projection_calculate_dispatch_timing(context, aggregate);
  fprintf(file,
          "    device=%u dispatches=%" PRIu64 " valid=%" PRIu64
          " invalid=%" PRIu64 " ticks[min/avg/max/total]=%" PRIu64
          "/%.3f/%" PRIu64 "/%" PRIu64,
          aggregate->physical_device_ordinal, aggregate->dispatch_count,
          aggregate->valid_count, aggregate->invalid_count,
          aggregate->valid_count ? aggregate->minimum_ticks : 0,
          timing.average_ticks,
          aggregate->valid_count ? aggregate->maximum_ticks : 0,
          aggregate->total_ticks);
  if (timing.has_time_ns) {
    fprintf(file,
            " ns[min/avg/max/total]=%" PRId64 "/%.3f/%" PRId64 "/%" PRId64,
            timing.minimum_ns, timing.average_ticks * timing.ns_per_tick,
            timing.maximum_ns, timing.total_ns);
  }
  fputc('\n', file);
}

static bool iree_profile_executable_print_text_dispatch_groups(
    const iree_profile_dispatch_context_t* context,
    const iree_profile_projection_index_t* index,
    iree_host_size_t function_index, FILE* file) {
  const iree_profile_projection_row_list_t* row_list =
      &index->functions[function_index].device_dispatches;
  bool has_aggregate = false;
  iree_host_size_t i = row_list->first_index;
  while (i != IREE_HOST_SIZE_MAX) {
    const iree_profile_dispatch_aggregate_t* aggregate =
        &context->aggregates[i];
    has_aggregate = true;
    iree_profile_executable_print_text_dispatch_group(context, aggregate, file);
    i = index->aggregate_next_indices[i];
  }
  return has_aggregate;
}

static void iree_profile_executable_print_text_host_dispatch_group(
    const iree_profile_host_dispatch_aggregate_t* aggregate, FILE* file) {
  fprintf(file,
          "    timing=host_execution device=%u dispatches=%" PRIu64
          " valid=%" PRIu64 " invalid=%" PRIu64
          " ns[min/avg/max/total]=%" PRId64 "/%.3f/%" PRId64 "/%" PRId64
          " tile_count=%" PRIu64 " tile_sum_ns=%" PRId64 "\n",
          aggregate->physical_device_ordinal, aggregate->dispatch_count,
          aggregate->valid_count, aggregate->invalid_count,
          aggregate->valid_count ? aggregate->minimum_ns : 0,
          aggregate->valid_count ? aggregate->mean_ns : 0.0,
          aggregate->valid_count ? aggregate->maximum_ns : 0,
          aggregate->total_ns, aggregate->total_tile_count,
          aggregate->total_tile_duration_sum_ns);
}

static bool iree_profile_executable_print_text_host_dispatch_groups(
    const iree_profile_dispatch_context_t* context,
    const iree_profile_projection_index_t* index,
    iree_host_size_t function_index, FILE* file) {
  const iree_profile_projection_row_list_t* row_list =
      &index->functions[function_index].host_dispatches;
  bool has_aggregate = false;
  iree_host_size_t i = row_list->first_index;
  while (i != IREE_HOST_SIZE_MAX) {
    const iree_profile_host_dispatch_aggregate_t* aggregate =
        &context->host_dispatch_aggregates[i];
    has_aggregate = true;
    iree_profile_executable_print_text_host_dispatch_group(aggregate, file);
    i = index->host_dispatch_next_indices[i];
  }
  return has_aggregate;
}

static void iree_profile_executable_print_text_functions(
    const iree_profile_dispatch_context_t* context,
    const iree_profile_projection_index_t* index,
    const iree_profile_model_executable_t* executable_info,
    iree_string_view_t filter, FILE* file) {
  iree_host_size_t i = executable_info->first_function_index;
  while (i != IREE_HOST_SIZE_MAX) {
    const iree_profile_model_function_t* function_info =
        &context->model.functions[i];
    char numeric_buffer[128];
    iree_string_view_t key = iree_profile_model_format_function_key(
        function_info, UINT32_MAX, numeric_buffer, sizeof(numeric_buffer));
    if (!iree_profile_key_matches(key, filter)) {
      i = function_info->next_function_index;
      continue;
    }

    iree_profile_executable_print_text_function(function_info, key, file);
    const bool has_device_dispatches =
        iree_profile_executable_print_text_dispatch_groups(context, index, i,
                                                           file);
    const bool has_host_dispatches =
        iree_profile_executable_print_text_host_dispatch_groups(context, index,
                                                                i, file);
    if (!has_device_dispatches && !has_host_dispatches) {
      fprintf(file, "    dispatches=0\n");
    }
    i = function_info->next_function_index;
  }
}

static iree_status_t iree_profile_executable_print_text(
    const iree_profile_dispatch_context_t* context, iree_string_view_t filter,
    int64_t id_filter, const iree_profile_projection_index_t* index,
    FILE* file) {
  iree_profile_executable_print_text_header(context, filter, file);

  for (iree_host_size_t i = 0; i < context->model.executable_count; ++i) {
    const iree_profile_model_executable_t* executable_info =
        &context->model.executables[i];
    const iree_hal_profile_executable_record_t* executable =
        &executable_info->record;
    if (!iree_profile_executable_matches_id(executable, id_filter)) {
      continue;
    }
    iree_profile_executable_print_text_executable(executable, file);
    iree_profile_executable_print_text_functions(context, index,
                                                 executable_info, filter, file);
  }
  return iree_ok_status();
}

static void iree_profile_executable_print_jsonl_summary(
    const iree_profile_dispatch_context_t* context, iree_string_view_t filter,
    FILE* file) {
  fprintf(file, "{\"type\":\"executable_summary\",\"filter\":");
  iree_profile_fprint_json_string(file, filter);
  fprintf(
      file,
      ",\"executables\":%" PRIhsz ",\"functions\":%" PRIhsz
      ",\"matched_dispatches\":%" PRIu64 ",\"valid_dispatches\":%" PRIu64
      ",\"invalid_dispatches\":%" PRIu64 ",\"matched_host_dispatches\":%" PRIu64
      ",\"valid_host_dispatches\":%" PRIu64
      ",\"invalid_host_dispatches\":%" PRIu64 "}\n",
      context->model.executable_count, context->model.function_count,
      context->matched_dispatch_count, context->valid_dispatch_count,
      context->invalid_dispatch_count, context->matched_host_dispatch_count,
      context->valid_host_dispatch_count, context->invalid_host_dispatch_count);
}

static void iree_profile_executable_print_jsonl_hash(const uint64_t hash[2],
                                                     bool has_hash,
                                                     FILE* file) {
  if (has_hash) {
    fputc('"', file);
    iree_profile_fprint_hash_hex(file, hash);
    fputc('"', file);
  } else {
    fprintf(file, "null");
  }
}

static void iree_profile_executable_print_jsonl_executable(
    const iree_hal_profile_executable_record_t* executable, FILE* file) {
  const bool has_code_object_hash = iree_all_bits_set(
      executable->flags, IREE_HAL_PROFILE_EXECUTABLE_FLAG_CODE_OBJECT_HASH);
  fprintf(file,
          "{\"type\":\"executable\",\"executable_id\":%" PRIu64
          ",\"flags\":%u,\"function_count\":%u"
          ",\"code_object_hash_present\":%s,\"code_object_hash\":",
          executable->executable_id, executable->flags,
          executable->function_count, has_code_object_hash ? "true" : "false");
  iree_profile_executable_print_jsonl_hash(executable->code_object_hash,
                                           has_code_object_hash, file);
  fputs("}\n", file);
}

static void iree_profile_executable_print_jsonl_function(
    const iree_profile_model_function_t* function_info, iree_string_view_t key,
    FILE* file) {
  const bool has_function_hash = iree_all_bits_set(
      function_info->flags,
      IREE_HAL_PROFILE_EXECUTABLE_FUNCTION_FLAG_FUNCTION_HASH);
  fprintf(file,
          "{\"type\":\"executable_function\",\"executable_id\":%" PRIu64
          ",\"function_ordinal\":%u,\"flags\":%u,\"key\":",
          function_info->executable_id, function_info->function_ordinal,
          function_info->flags);
  iree_profile_fprint_json_string(file, key);
  fprintf(file,
          ",\"constant_count\":%u,\"binding_count\":%u"
          ",\"parameter_count\":%u,\"workgroup_size\":[%u,%u,%u]"
          ",\"function_hash_present\":%s,\"function_hash\":",
          function_info->constant_count, function_info->binding_count,
          function_info->parameter_count, function_info->workgroup_size[0],
          function_info->workgroup_size[1], function_info->workgroup_size[2],
          has_function_hash ? "true" : "false");
  iree_profile_executable_print_jsonl_hash(function_info->function_hash,
                                           has_function_hash, file);
  fputs("}\n", file);
}

static void iree_profile_executable_print_jsonl_dispatch_group(
    const iree_profile_dispatch_context_t* context,
    const iree_profile_dispatch_aggregate_t* aggregate, iree_string_view_t key,
    FILE* file) {
  const iree_profile_projection_dispatch_timing_t timing =
      iree_profile_projection_calculate_dispatch_timing(context, aggregate);
  fprintf(file,
          "{\"type\":\"executable_function_dispatch_group\""
          ",\"physical_device_ordinal\":%u"
          ",\"executable_id\":%" PRIu64
          ",\"function_ordinal\":%u"
          ",\"key\":",
          aggregate->physical_device_ordinal, aggregate->executable_id,
          aggregate->function_ordinal);
  iree_profile_fprint_json_string(file, key);
  fprintf(file,
          ",\"dispatches\":%" PRIu64 ",\"valid\":%" PRIu64
          ",\"invalid\":%" PRIu64 ",\"min_ticks\":%" PRIu64
          ",\"avg_ticks\":%.3f,\"max_ticks\":%" PRIu64
          ",\"total_ticks\":%" PRIu64
          ",\"clock_fit_available\":%s"
          ",\"time_ns_available\":%s"
          ",\"min_ns\":%" PRId64
          ",\"avg_ns\":%.3f"
          ",\"max_ns\":%" PRId64 ",\"total_ns\":%" PRId64 "}\n",
          aggregate->dispatch_count, aggregate->valid_count,
          aggregate->invalid_count,
          aggregate->valid_count ? aggregate->minimum_ticks : 0,
          timing.average_ticks,
          aggregate->valid_count ? aggregate->maximum_ticks : 0,
          aggregate->total_ticks, timing.has_clock_fit ? "true" : "false",
          timing.has_time_ns ? "true" : "false",
          timing.has_time_ns ? timing.minimum_ns : 0,
          timing.has_time_ns ? timing.average_ticks * timing.ns_per_tick : 0.0,
          timing.has_time_ns ? timing.maximum_ns : 0,
          timing.has_time_ns ? timing.total_ns : 0);
}

static void iree_profile_executable_print_jsonl_dispatch_groups(
    const iree_profile_dispatch_context_t* context,
    const iree_profile_projection_index_t* index,
    iree_host_size_t function_index, iree_string_view_t key, FILE* file) {
  const iree_profile_projection_row_list_t* row_list =
      &index->functions[function_index].device_dispatches;
  iree_host_size_t i = row_list->first_index;
  while (i != IREE_HOST_SIZE_MAX) {
    const iree_profile_dispatch_aggregate_t* aggregate =
        &context->aggregates[i];
    iree_profile_executable_print_jsonl_dispatch_group(context, aggregate, key,
                                                       file);
    i = index->aggregate_next_indices[i];
  }
}

static void iree_profile_executable_print_jsonl_host_dispatch_group(
    const iree_profile_host_dispatch_aggregate_t* aggregate,
    iree_string_view_t key, FILE* file) {
  const double stddev_ns =
      iree_profile_projection_host_dispatch_stddev_ns(aggregate);
  fprintf(file,
          "{\"type\":\"executable_function_host_dispatch_group\""
          ",\"timing_source\":\"host_execution_event\""
          ",\"time_domain\":\"iree_host_time_ns\""
          ",\"duration_time_domain\":\"iree_host_duration_ns\""
          ",\"physical_device_ordinal\":%u"
          ",\"executable_id\":%" PRIu64
          ",\"function_ordinal\":%u"
          ",\"key\":",
          aggregate->physical_device_ordinal, aggregate->executable_id,
          aggregate->function_ordinal);
  iree_profile_fprint_json_string(file, key);
  fprintf(file,
          ",\"dispatches\":%" PRIu64 ",\"valid\":%" PRIu64
          ",\"invalid\":%" PRIu64 ",\"min_ns\":%" PRId64
          ",\"avg_ns\":%.3f,\"stddev_ns\":%.3f"
          ",\"max_ns\":%" PRId64 ",\"total_ns\":%" PRId64
          ",\"total_tile_count\":%" PRIu64
          ",\"total_tile_duration_sum_ns\":%" PRId64 "}\n",
          aggregate->dispatch_count, aggregate->valid_count,
          aggregate->invalid_count,
          aggregate->valid_count ? aggregate->minimum_ns : 0,
          aggregate->valid_count ? aggregate->mean_ns : 0.0, stddev_ns,
          aggregate->valid_count ? aggregate->maximum_ns : 0,
          aggregate->total_ns, aggregate->total_tile_count,
          aggregate->total_tile_duration_sum_ns);
}

static void iree_profile_executable_print_jsonl_host_dispatch_groups(
    const iree_profile_dispatch_context_t* context,
    const iree_profile_projection_index_t* index,
    iree_host_size_t function_index, iree_string_view_t key, FILE* file) {
  const iree_profile_projection_row_list_t* row_list =
      &index->functions[function_index].host_dispatches;
  iree_host_size_t i = row_list->first_index;
  while (i != IREE_HOST_SIZE_MAX) {
    const iree_profile_host_dispatch_aggregate_t* aggregate =
        &context->host_dispatch_aggregates[i];
    iree_profile_executable_print_jsonl_host_dispatch_group(aggregate, key,
                                                            file);
    i = index->host_dispatch_next_indices[i];
  }
}

static void iree_profile_executable_print_jsonl_functions(
    const iree_profile_dispatch_context_t* context,
    const iree_profile_projection_index_t* index,
    const iree_profile_model_executable_t* executable_info,
    iree_string_view_t filter, FILE* file) {
  iree_host_size_t i = executable_info->first_function_index;
  while (i != IREE_HOST_SIZE_MAX) {
    const iree_profile_model_function_t* function_info =
        &context->model.functions[i];
    char numeric_buffer[128];
    iree_string_view_t key = iree_profile_model_format_function_key(
        function_info, UINT32_MAX, numeric_buffer, sizeof(numeric_buffer));
    if (!iree_profile_key_matches(key, filter)) {
      i = function_info->next_function_index;
      continue;
    }

    iree_profile_executable_print_jsonl_function(function_info, key, file);
    iree_profile_executable_print_jsonl_dispatch_groups(context, index, i, key,
                                                        file);
    iree_profile_executable_print_jsonl_host_dispatch_groups(context, index, i,
                                                             key, file);
    i = function_info->next_function_index;
  }
}

static iree_status_t iree_profile_executable_print_jsonl(
    const iree_profile_dispatch_context_t* context, iree_string_view_t filter,
    int64_t id_filter, const iree_profile_projection_index_t* index,
    FILE* file) {
  iree_profile_executable_print_jsonl_summary(context, filter, file);
  for (iree_host_size_t i = 0; i < context->model.executable_count; ++i) {
    const iree_profile_model_executable_t* executable_info =
        &context->model.executables[i];
    const iree_hal_profile_executable_record_t* executable =
        &executable_info->record;
    if (!iree_profile_executable_matches_id(executable, id_filter)) {
      continue;
    }
    iree_profile_executable_print_jsonl_executable(executable, file);
    iree_profile_executable_print_jsonl_functions(
        context, index, executable_info, filter, file);
  }
  return iree_ok_status();
}

static bool iree_profile_command_buffer_matches_id(
    const iree_hal_profile_command_buffer_record_t* command_buffer,
    int64_t id_filter) {
  return id_filter < 0 ||
         command_buffer->command_buffer_id == (uint64_t)id_filter;
}

static void iree_profile_command_print_text_header(
    const iree_profile_dispatch_context_t* context, iree_string_view_t filter,
    iree_host_size_t matched_command_operation_count, FILE* file) {
  fprintf(file, "IREE HAL profile command-buffer summary\n");
  fprintf(file, "filter: %.*s\n", (int)filter.size, filter.data);
  fprintf(
      file,
      "command_buffers=%" PRIhsz " executions=%" PRIhsz
      " host_executions=%" PRIhsz " command_operations=%" PRIhsz
      " matched_command_operations=%" PRIhsz " matched_dispatches=%" PRIu64
      " valid=%" PRIu64 " invalid=%" PRIu64 " matched_host_dispatches=%" PRIu64
      " host_valid=%" PRIu64 " host_invalid=%" PRIu64 "\n",
      context->model.command_buffer_count, context->command_aggregate_count,
      context->host_command_aggregate_count,
      context->model.command_operation_count, matched_command_operation_count,
      context->matched_dispatch_count, context->valid_dispatch_count,
      context->invalid_dispatch_count, context->matched_host_dispatch_count,
      context->valid_host_dispatch_count, context->invalid_host_dispatch_count);
}

static void iree_profile_command_print_text_command_buffer(
    const iree_hal_profile_command_buffer_record_t* command_buffer,
    FILE* file) {
  fprintf(file,
          "command_buffer %" PRIu64 ": device=%u mode=%" PRIu64
          " categories=%" PRIu64 " queue_affinity=%" PRIu64 "\n",
          command_buffer->command_buffer_id,
          command_buffer->physical_device_ordinal, command_buffer->mode,
          command_buffer->command_categories, command_buffer->queue_affinity);
}

static iree_status_t iree_profile_command_print_text_operations(
    const iree_profile_dispatch_context_t* context,
    const iree_hal_profile_command_buffer_record_t* command_buffer,
    iree_string_view_t filter, FILE* file) {
  const iree_profile_model_command_buffer_t* command_buffer_info =
      iree_profile_model_find_command_buffer(&context->model,
                                             command_buffer->command_buffer_id);
  if (!command_buffer_info) {
    return iree_make_status(
        IREE_STATUS_DATA_LOSS,
        "command-buffer report references missing command-buffer metadata "
        "command_buffer=%" PRIu64,
        command_buffer->command_buffer_id);
  }

  iree_status_t status = iree_ok_status();
  iree_host_size_t operation_index = command_buffer_info->first_operation_index;
  while (operation_index != IREE_HOST_SIZE_MAX && iree_status_is_ok(status)) {
    const iree_hal_profile_command_operation_record_t* operation =
        &context->model.command_operations[operation_index].record;
    status = iree_profile_command_print_operation_text(context, operation,
                                                       filter, file);
    operation_index =
        context->model.command_operations[operation_index].next_operation_index;
  }
  return status;
}

static void iree_profile_command_print_text_execution(
    const iree_profile_dispatch_context_t* context,
    const iree_profile_dispatch_command_aggregate_t* aggregate, FILE* file) {
  const iree_profile_projection_submission_timing_t timing =
      iree_profile_projection_calculate_submission_timing(
          context, aggregate->physical_device_ordinal,
          aggregate->earliest_start_tick, aggregate->latest_end_tick,
          aggregate->valid_count, aggregate->total_ticks);
  fprintf(file,
          "  submission=%" PRIu64 " queue=%u stream=%" PRIu64
          " dispatches=%" PRIu64 " valid=%" PRIu64 " invalid=%" PRIu64
          " span_ticks=%.3f total_dispatch_ticks=%" PRIu64,
          aggregate->submission_id, aggregate->queue_ordinal,
          aggregate->stream_id, aggregate->dispatch_count,
          aggregate->valid_count, aggregate->invalid_count, timing.span_ticks,
          aggregate->total_ticks);
  if (timing.has_total_dispatch_ns) {
    fprintf(file, " span_ns=%.3f total_dispatch_ns=%" PRId64,
            timing.span_ticks * timing.ns_per_tick, timing.total_dispatch_ns);
  }
  fputc('\n', file);
}

static void iree_profile_command_print_text_executions(
    const iree_profile_dispatch_context_t* context,
    const iree_profile_projection_index_t* index,
    iree_host_size_t command_buffer_index, FILE* file) {
  const iree_profile_projection_row_list_t* row_list =
      &index->command_buffers[command_buffer_index].device_executions;
  iree_host_size_t i = row_list->first_index;
  while (i != IREE_HOST_SIZE_MAX) {
    const iree_profile_dispatch_command_aggregate_t* aggregate =
        &context->command_aggregates[i];
    iree_profile_command_print_text_execution(context, aggregate, file);
    i = index->command_next_indices[i];
  }
}

static void iree_profile_command_print_text_host_execution(
    const iree_profile_host_dispatch_command_aggregate_t* aggregate,
    FILE* file) {
  const int64_t span_ns = aggregate->valid_count != 0
                              ? aggregate->latest_end_host_time_ns -
                                    aggregate->earliest_start_host_time_ns
                              : 0;
  fprintf(file,
          "  host_submission=%" PRIu64 " queue=%u stream=%" PRIu64
          " dispatches=%" PRIu64 " valid=%" PRIu64 " invalid=%" PRIu64
          " span_ns=%" PRId64 " total_dispatch_ns=%" PRId64
          " tile_count=%" PRIu64 " tile_sum_ns=%" PRId64 "\n",
          aggregate->submission_id, aggregate->queue_ordinal,
          aggregate->stream_id, aggregate->dispatch_count,
          aggregate->valid_count, aggregate->invalid_count, span_ns,
          aggregate->total_ns, aggregate->total_tile_count,
          aggregate->total_tile_duration_sum_ns);
}

static void iree_profile_command_print_text_host_executions(
    const iree_profile_dispatch_context_t* context,
    const iree_profile_projection_index_t* index,
    iree_host_size_t command_buffer_index, FILE* file) {
  const iree_profile_projection_row_list_t* row_list =
      &index->command_buffers[command_buffer_index].host_executions;
  iree_host_size_t i = row_list->first_index;
  while (i != IREE_HOST_SIZE_MAX) {
    const iree_profile_host_dispatch_command_aggregate_t* aggregate =
        &context->host_command_aggregates[i];
    iree_profile_command_print_text_host_execution(aggregate, file);
    i = index->host_command_next_indices[i];
  }
}

static iree_status_t iree_profile_command_print_text(
    const iree_profile_dispatch_context_t* context, iree_string_view_t filter,
    int64_t id_filter, const iree_profile_projection_index_t* index,
    FILE* file) {
  iree_host_size_t matched_command_operation_count = 0;
  IREE_RETURN_IF_ERROR(iree_profile_command_count_matching_operations(
      context, filter, id_filter, &matched_command_operation_count));

  iree_profile_command_print_text_header(context, filter,
                                         matched_command_operation_count, file);

  iree_status_t status = iree_ok_status();
  for (iree_host_size_t i = 0;
       i < context->model.command_buffer_count && iree_status_is_ok(status);
       ++i) {
    const iree_hal_profile_command_buffer_record_t* command_buffer =
        &context->model.command_buffers[i].record;
    if (!iree_profile_command_buffer_matches_id(command_buffer, id_filter)) {
      continue;
    }
    iree_profile_command_print_text_command_buffer(command_buffer, file);
    status = iree_profile_command_print_text_operations(context, command_buffer,
                                                        filter, file);
    if (iree_status_is_ok(status)) {
      iree_profile_command_print_text_executions(context, index, i, file);
      iree_profile_command_print_text_host_executions(context, index, i, file);
    }
  }
  return status;
}

static void iree_profile_command_print_jsonl_summary(
    const iree_profile_dispatch_context_t* context, iree_string_view_t filter,
    iree_host_size_t matched_command_operation_count, FILE* file) {
  fprintf(file, "{\"type\":\"command_summary\",\"filter\":");
  iree_profile_fprint_json_string(file, filter);
  fprintf(
      file,
      ",\"command_buffers\":%" PRIhsz ",\"executions\":%" PRIhsz
      ",\"host_executions\":%" PRIhsz ",\"command_operations\":%" PRIhsz
      ",\"matched_command_operations\":%" PRIhsz
      ",\"matched_dispatches\":%" PRIu64 ",\"valid_dispatches\":%" PRIu64
      ",\"invalid_dispatches\":%" PRIu64 ",\"matched_host_dispatches\":%" PRIu64
      ",\"valid_host_dispatches\":%" PRIu64
      ",\"invalid_host_dispatches\":%" PRIu64 "}\n",
      context->model.command_buffer_count, context->command_aggregate_count,
      context->host_command_aggregate_count,
      context->model.command_operation_count, matched_command_operation_count,
      context->matched_dispatch_count, context->valid_dispatch_count,
      context->invalid_dispatch_count, context->matched_host_dispatch_count,
      context->valid_host_dispatch_count, context->invalid_host_dispatch_count);
}

static void iree_profile_command_print_jsonl_command_buffer(
    const iree_hal_profile_command_buffer_record_t* command_buffer,
    FILE* file) {
  fprintf(file,
          "{\"type\":\"command_buffer\",\"command_buffer_id\":%" PRIu64
          ",\"physical_device_ordinal\":%u,\"mode\":%" PRIu64
          ",\"command_categories\":%" PRIu64 ",\"queue_affinity\":%" PRIu64
          "}\n",
          command_buffer->command_buffer_id,
          command_buffer->physical_device_ordinal, command_buffer->mode,
          command_buffer->command_categories, command_buffer->queue_affinity);
}

static void iree_profile_command_print_jsonl_command_buffers(
    const iree_profile_dispatch_context_t* context, int64_t id_filter,
    FILE* file) {
  for (iree_host_size_t i = 0; i < context->model.command_buffer_count; ++i) {
    const iree_hal_profile_command_buffer_record_t* command_buffer =
        &context->model.command_buffers[i].record;
    if (!iree_profile_command_buffer_matches_id(command_buffer, id_filter)) {
      continue;
    }
    iree_profile_command_print_jsonl_command_buffer(command_buffer, file);
  }
}

static iree_status_t iree_profile_command_print_jsonl_operations(
    const iree_profile_dispatch_context_t* context, iree_string_view_t filter,
    int64_t id_filter, FILE* file) {
  iree_status_t status = iree_ok_status();

  if (id_filter >= 0) {
    const iree_profile_model_command_buffer_t* command_buffer =
        iree_profile_model_find_command_buffer(&context->model,
                                               (uint64_t)id_filter);
    iree_host_size_t operation_index =
        command_buffer ? command_buffer->first_operation_index
                       : IREE_HOST_SIZE_MAX;
    while (operation_index != IREE_HOST_SIZE_MAX && iree_status_is_ok(status)) {
      const iree_hal_profile_command_operation_record_t* operation =
          &context->model.command_operations[operation_index].record;
      status = iree_profile_command_print_operation_jsonl(context, operation,
                                                          filter, file);
      operation_index = context->model.command_operations[operation_index]
                            .next_operation_index;
    }
    return status;
  }

  for (iree_host_size_t i = 0;
       i < context->model.command_operation_count && iree_status_is_ok(status);
       ++i) {
    const iree_hal_profile_command_operation_record_t* operation =
        &context->model.command_operations[i].record;
    status = iree_profile_command_print_operation_jsonl(context, operation,
                                                        filter, file);
  }
  return status;
}

static void iree_profile_command_print_jsonl_execution(
    const iree_profile_dispatch_context_t* context,
    const iree_profile_dispatch_command_aggregate_t* aggregate, FILE* file) {
  const iree_profile_projection_submission_timing_t timing =
      iree_profile_projection_calculate_submission_timing(
          context, aggregate->physical_device_ordinal,
          aggregate->earliest_start_tick, aggregate->latest_end_tick,
          aggregate->valid_count, aggregate->total_ticks);
  fprintf(file,
          "{\"type\":\"command_execution\",\"command_buffer_id\":%" PRIu64
          ",\"submission_id\":%" PRIu64
          ",\"physical_device_ordinal\":%u"
          ",\"queue_ordinal\":%u,\"stream_id\":%" PRIu64
          ",\"dispatches\":%" PRIu64 ",\"valid\":%" PRIu64
          ",\"invalid\":%" PRIu64
          ",\"span_ticks\":%.3f"
          ",\"total_dispatch_ticks\":%" PRIu64
          ",\"clock_fit_available\":%s"
          ",\"total_dispatch_time_ns_available\":%s"
          ",\"span_ns\":%.3f,\"total_dispatch_ns\":%" PRId64 "}\n",
          aggregate->command_buffer_id, aggregate->submission_id,
          aggregate->physical_device_ordinal, aggregate->queue_ordinal,
          aggregate->stream_id, aggregate->dispatch_count,
          aggregate->valid_count, aggregate->invalid_count, timing.span_ticks,
          aggregate->total_ticks, timing.has_clock_fit ? "true" : "false",
          timing.has_total_dispatch_ns ? "true" : "false",
          timing.has_clock_fit ? timing.span_ticks * timing.ns_per_tick : 0.0,
          timing.has_total_dispatch_ns ? timing.total_dispatch_ns : 0);
}

static void iree_profile_command_print_jsonl_executions(
    const iree_profile_dispatch_context_t* context,
    const iree_profile_projection_index_t* index, int64_t id_filter,
    FILE* file) {
  if (id_filter >= 0) {
    const iree_profile_model_command_buffer_t* command_buffer =
        iree_profile_model_find_command_buffer(&context->model,
                                               (uint64_t)id_filter);
    if (!command_buffer) return;
    const iree_host_size_t command_buffer_index =
        (iree_host_size_t)(command_buffer - context->model.command_buffers);
    const iree_profile_projection_row_list_t* row_list =
        &index->command_buffers[command_buffer_index].device_executions;
    iree_host_size_t i = row_list->first_index;
    while (i != IREE_HOST_SIZE_MAX) {
      const iree_profile_dispatch_command_aggregate_t* aggregate =
          &context->command_aggregates[i];
      iree_profile_command_print_jsonl_execution(context, aggregate, file);
      i = index->command_next_indices[i];
    }
    return;
  }

  for (iree_host_size_t i = 0; i < context->command_aggregate_count; ++i) {
    const iree_profile_dispatch_command_aggregate_t* aggregate =
        &context->command_aggregates[i];
    iree_profile_command_print_jsonl_execution(context, aggregate, file);
  }
}

static void iree_profile_command_print_jsonl_host_execution(
    const iree_profile_host_dispatch_command_aggregate_t* aggregate,
    FILE* file) {
  const int64_t span_ns = aggregate->valid_count != 0
                              ? aggregate->latest_end_host_time_ns -
                                    aggregate->earliest_start_host_time_ns
                              : 0;
  fprintf(
      file,
      "{\"type\":\"command_host_execution\""
      ",\"timing_source\":\"host_execution_event\""
      ",\"time_domain\":\"iree_host_time_ns\""
      ",\"duration_time_domain\":\"iree_host_duration_ns\""
      ",\"command_buffer_id\":%" PRIu64 ",\"submission_id\":%" PRIu64
      ",\"physical_device_ordinal\":%u"
      ",\"queue_ordinal\":%u,\"stream_id\":%" PRIu64 ",\"dispatches\":%" PRIu64
      ",\"valid\":%" PRIu64 ",\"invalid\":%" PRIu64 ",\"span_ns\":%" PRId64
      ",\"total_dispatch_ns\":%" PRId64 ",\"total_tile_count\":%" PRIu64
      ",\"total_tile_duration_sum_ns\":%" PRId64 "}\n",
      aggregate->command_buffer_id, aggregate->submission_id,
      aggregate->physical_device_ordinal, aggregate->queue_ordinal,
      aggregate->stream_id, aggregate->dispatch_count, aggregate->valid_count,
      aggregate->invalid_count, span_ns, aggregate->total_ns,
      aggregate->total_tile_count, aggregate->total_tile_duration_sum_ns);
}

static void iree_profile_command_print_jsonl_host_executions(
    const iree_profile_dispatch_context_t* context,
    const iree_profile_projection_index_t* index, int64_t id_filter,
    FILE* file) {
  if (id_filter >= 0) {
    const iree_profile_model_command_buffer_t* command_buffer =
        iree_profile_model_find_command_buffer(&context->model,
                                               (uint64_t)id_filter);
    if (!command_buffer) return;
    const iree_host_size_t command_buffer_index =
        (iree_host_size_t)(command_buffer - context->model.command_buffers);
    const iree_profile_projection_row_list_t* row_list =
        &index->command_buffers[command_buffer_index].host_executions;
    iree_host_size_t i = row_list->first_index;
    while (i != IREE_HOST_SIZE_MAX) {
      const iree_profile_host_dispatch_command_aggregate_t* aggregate =
          &context->host_command_aggregates[i];
      iree_profile_command_print_jsonl_host_execution(aggregate, file);
      i = index->host_command_next_indices[i];
    }
    return;
  }

  for (iree_host_size_t i = 0; i < context->host_command_aggregate_count; ++i) {
    const iree_profile_host_dispatch_command_aggregate_t* aggregate =
        &context->host_command_aggregates[i];
    iree_profile_command_print_jsonl_host_execution(aggregate, file);
  }
}

static iree_status_t iree_profile_command_print_jsonl(
    const iree_profile_dispatch_context_t* context, iree_string_view_t filter,
    int64_t id_filter, const iree_profile_projection_index_t* index,
    FILE* file) {
  iree_host_size_t matched_command_operation_count = 0;
  IREE_RETURN_IF_ERROR(iree_profile_command_count_matching_operations(
      context, filter, id_filter, &matched_command_operation_count));

  iree_profile_command_print_jsonl_summary(
      context, filter, matched_command_operation_count, file);
  iree_profile_command_print_jsonl_command_buffers(context, id_filter, file);

  iree_status_t status = iree_profile_command_print_jsonl_operations(
      context, filter, id_filter, file);
  if (iree_status_is_ok(status)) {
    iree_profile_command_print_jsonl_executions(context, index, id_filter,
                                                file);
    iree_profile_command_print_jsonl_host_executions(context, index, id_filter,
                                                     file);
  }
  return status;
}

static void iree_profile_queue_print_text_header(
    const iree_profile_dispatch_context_t* context, iree_string_view_t filter,
    FILE* file) {
  fprintf(file, "IREE HAL profile queue summary\n");
  fprintf(file, "filter: %.*s\n", (int)filter.size, filter.data);
  fprintf(file,
          "queues=%" PRIhsz " queue_events=%" PRIhsz
          " queue_device_events=%" PRIhsz " host_execution_events=%" PRIhsz
          " truncated_chunks=%" PRIu64 " dropped_records=%" PRIu64
          " submissions=%" PRIhsz " matched_dispatches=%" PRIu64
          " valid=%" PRIu64 " invalid=%" PRIu64 "\n",
          context->model.queue_count, context->queue_query.queue_event_count,
          context->queue_query.queue_device_event_count,
          context->queue_query.host_execution_event_count,
          context->queue_query.truncated_chunk_count,
          context->queue_query.dropped_record_count,
          context->queue_aggregate_count, context->matched_dispatch_count,
          context->valid_dispatch_count, context->invalid_dispatch_count);
  fprintf(file,
          "host_dispatches: matched=%" PRIu64 " valid=%" PRIu64
          " invalid=%" PRIu64 " command_executions=%" PRIhsz "\n",
          context->matched_host_dispatch_count,
          context->valid_host_dispatch_count,
          context->invalid_host_dispatch_count,
          context->host_command_aggregate_count);
}

static void iree_profile_queue_print_text_submission(
    const iree_profile_dispatch_context_t* context,
    const iree_profile_dispatch_queue_aggregate_t* aggregate, FILE* file) {
  const iree_profile_projection_submission_timing_t timing =
      iree_profile_projection_calculate_submission_timing(
          context, aggregate->physical_device_ordinal,
          aggregate->earliest_start_tick, aggregate->latest_end_tick,
          aggregate->valid_count, aggregate->total_ticks);
  fprintf(file,
          "  submission=%" PRIu64 " dispatches=%" PRIu64 " valid=%" PRIu64
          " invalid=%" PRIu64 " span_ticks=%.3f total_dispatch_ticks=%" PRIu64,
          aggregate->submission_id, aggregate->dispatch_count,
          aggregate->valid_count, aggregate->invalid_count, timing.span_ticks,
          aggregate->total_ticks);
  if (timing.has_total_dispatch_ns) {
    fprintf(file, " span_ns=%.3f total_dispatch_ns=%" PRId64,
            timing.span_ticks * timing.ns_per_tick, timing.total_dispatch_ns);
  }
  fputc('\n', file);
}

static void iree_profile_queue_print_text_device_event(
    const iree_profile_dispatch_context_t* context,
    const iree_hal_profile_queue_device_event_t* event, FILE* file) {
  const bool is_valid = event->start_tick != 0 && event->end_tick != 0 &&
                        event->end_tick >= event->start_tick;
  const uint64_t duration_ticks =
      is_valid ? event->end_tick - event->start_tick : 0;
  iree_profile_model_clock_fit_t clock_fit;
  const bool has_clock_fit =
      iree_profile_projection_try_fit_driver_host_cpu_clock(
          context, event->physical_device_ordinal, &clock_fit);
  int64_t duration_ns = 0;
  const bool has_duration_ns = is_valid && has_clock_fit &&
                               iree_profile_model_clock_fit_scale_ticks_to_ns(
                                   &clock_fit, duration_ticks, &duration_ns);
  fprintf(file,
          "  device_event=%" PRIu64 " type=%s submission=%" PRIu64
          " command_buffer=%" PRIu64 " allocation=%" PRIu64
          " ops=%u ticks=%" PRIu64 " valid=%s",
          event->event_id, iree_profile_queue_event_type_name(event->type),
          event->submission_id, event->command_buffer_id, event->allocation_id,
          event->operation_count, duration_ticks, is_valid ? "true" : "false");
  if (has_duration_ns) {
    fprintf(file, " ns=%" PRId64, duration_ns);
  }
  fputc('\n', file);
}

static void iree_profile_queue_print_text_host_execution_event(
    const iree_hal_profile_host_execution_event_t* event, FILE* file) {
  const bool is_valid = event->start_host_time_ns >= 0 &&
                        event->end_host_time_ns >= event->start_host_time_ns;
  const int64_t duration_ns =
      is_valid ? event->end_host_time_ns - event->start_host_time_ns : 0;
  fprintf(file,
          "  host_execution_event=%" PRIu64 " type=%s submission=%" PRIu64
          " command_buffer=%" PRIu64 " allocation=%" PRIu64
          " ops=%u tiles=%" PRIu64 " ns=%" PRId64 " tile_sum_ns=%" PRId64
          " valid=%s\n",
          event->event_id, iree_profile_queue_event_type_name(event->type),
          event->submission_id, event->command_buffer_id, event->allocation_id,
          event->operation_count, event->tile_count, duration_ns,
          event->tile_duration_sum_ns, is_valid ? "true" : "false");
}

static void iree_profile_queue_print_text_event(
    const iree_hal_profile_queue_event_t* event, FILE* file) {
  fprintf(
      file,
      "  event=%" PRIu64 " type=%s submission=%" PRIu64
      " strategy=%s submit=%" PRId64 " ready=%" PRId64
      " waits=%u signals=%u barriers=%u ops=%u bytes=%" PRIu64 "\n",
      event->event_id, iree_profile_queue_event_type_name(event->type),
      event->submission_id,
      iree_profile_queue_dependency_strategy_name(event->dependency_strategy),
      event->host_time_ns, event->ready_host_time_ns, event->wait_count,
      event->signal_count, event->barrier_count, event->operation_count,
      event->payload_length);
}

static void iree_profile_queue_print_text_submissions_for_queue(
    const iree_profile_dispatch_context_t* context,
    const iree_profile_projection_index_t* index, iree_host_size_t queue_index,
    FILE* file) {
  const iree_profile_projection_row_list_t* row_list =
      &index->queues[queue_index].submissions;
  iree_host_size_t i = row_list->first_index;
  while (i != IREE_HOST_SIZE_MAX) {
    const iree_profile_dispatch_queue_aggregate_t* aggregate =
        &context->queue_aggregates[i];
    iree_profile_queue_print_text_submission(context, aggregate, file);
    i = index->queue_next_indices[i];
  }
}

static void iree_profile_queue_print_text_device_events_for_queue(
    const iree_profile_dispatch_context_t* context,
    const iree_profile_projection_index_t* index, iree_host_size_t queue_index,
    FILE* file) {
  const iree_profile_projection_row_list_t* row_list =
      &index->queues[queue_index].device_events;
  iree_host_size_t i = row_list->first_index;
  while (i != IREE_HOST_SIZE_MAX) {
    const iree_hal_profile_queue_device_event_t* event =
        &context->queue_query.queue_device_events[i].record;
    iree_profile_queue_print_text_device_event(context, event, file);
    i = index->queue_device_event_next_indices[i];
  }
}

static void iree_profile_queue_print_text_host_execution_events_for_queue(
    const iree_profile_dispatch_context_t* context,
    const iree_profile_projection_index_t* index, iree_host_size_t queue_index,
    FILE* file) {
  const iree_profile_projection_row_list_t* row_list =
      &index->queues[queue_index].host_execution_events;
  iree_host_size_t i = row_list->first_index;
  while (i != IREE_HOST_SIZE_MAX) {
    const iree_hal_profile_host_execution_event_t* event =
        &context->queue_query.host_execution_events[i].record;
    iree_profile_queue_print_text_host_execution_event(event, file);
    i = index->host_execution_event_next_indices[i];
  }
}

static void iree_profile_queue_print_text_events_for_queue(
    const iree_profile_dispatch_context_t* context,
    const iree_profile_projection_index_t* index, iree_host_size_t queue_index,
    FILE* file) {
  const iree_profile_projection_row_list_t* row_list =
      &index->queues[queue_index].events;
  iree_host_size_t i = row_list->first_index;
  while (i != IREE_HOST_SIZE_MAX) {
    const iree_hal_profile_queue_event_t* event =
        &context->queue_query.queue_events[i].record;
    iree_profile_queue_print_text_event(event, file);
    i = index->queue_event_next_indices[i];
  }
}

static void iree_profile_queue_print_text_queue(
    const iree_profile_dispatch_context_t* context,
    const iree_profile_projection_index_t* index, iree_host_size_t queue_index,
    FILE* file) {
  const iree_hal_profile_queue_record_t* queue =
      &context->model.queues[queue_index].record;
  fprintf(file, "queue device=%u ordinal=%u stream=%" PRIu64 "\n",
          queue->physical_device_ordinal, queue->queue_ordinal,
          queue->stream_id);

  iree_profile_queue_print_text_submissions_for_queue(context, index,
                                                      queue_index, file);
  iree_profile_queue_print_text_device_events_for_queue(context, index,
                                                        queue_index, file);
  iree_profile_queue_print_text_host_execution_events_for_queue(
      context, index, queue_index, file);
  iree_profile_queue_print_text_events_for_queue(context, index, queue_index,
                                                 file);
}

static iree_status_t iree_profile_queue_print_text(
    const iree_profile_dispatch_context_t* context, iree_string_view_t filter,
    const iree_profile_projection_index_t* index, FILE* file) {
  iree_profile_queue_print_text_header(context, filter, file);
  for (iree_host_size_t i = 0; i < context->model.queue_count; ++i) {
    iree_profile_queue_print_text_queue(context, index, i, file);
  }
  return iree_ok_status();
}

static void iree_profile_queue_print_jsonl_summary(
    const iree_profile_dispatch_context_t* context, iree_string_view_t filter,
    FILE* file) {
  fprintf(file, "{\"type\":\"queue_summary\",\"filter\":");
  iree_profile_fprint_json_string(file, filter);
  fprintf(
      file,
      ",\"queues\":%" PRIhsz ",\"queue_events\":%" PRIhsz
      ",\"queue_device_events\":%" PRIhsz ",\"host_execution_events\":%" PRIhsz
      ",\"truncated_chunks\":%" PRIu64 ",\"dropped_records\":%" PRIu64
      ",\"submissions\":%" PRIhsz ",\"matched_dispatches\":%" PRIu64
      ",\"valid_dispatches\":%" PRIu64 ",\"invalid_dispatches\":%" PRIu64
      ",\"matched_host_dispatches\":%" PRIu64
      ",\"valid_host_dispatches\":%" PRIu64
      ",\"invalid_host_dispatches\":%" PRIu64
      ",\"host_command_executions\":%" PRIhsz "}\n",
      context->model.queue_count, context->queue_query.queue_event_count,
      context->queue_query.queue_device_event_count,
      context->queue_query.host_execution_event_count,
      context->queue_query.truncated_chunk_count,
      context->queue_query.dropped_record_count, context->queue_aggregate_count,
      context->matched_dispatch_count, context->valid_dispatch_count,
      context->invalid_dispatch_count, context->matched_host_dispatch_count,
      context->valid_host_dispatch_count, context->invalid_host_dispatch_count,
      context->host_command_aggregate_count);
}

static void iree_profile_queue_print_jsonl_queues(
    const iree_profile_dispatch_context_t* context, FILE* file) {
  for (iree_host_size_t i = 0; i < context->model.queue_count; ++i) {
    const iree_hal_profile_queue_record_t* queue =
        &context->model.queues[i].record;
    fprintf(file,
            "{\"type\":\"queue\",\"physical_device_ordinal\":%u"
            ",\"queue_ordinal\":%u,\"stream_id\":%" PRIu64 "}\n",
            queue->physical_device_ordinal, queue->queue_ordinal,
            queue->stream_id);
  }
}

static void iree_profile_queue_print_jsonl_submissions(
    const iree_profile_dispatch_context_t* context, int64_t id_filter,
    FILE* file) {
  for (iree_host_size_t i = 0; i < context->queue_aggregate_count; ++i) {
    const iree_profile_dispatch_queue_aggregate_t* aggregate =
        &context->queue_aggregates[i];
    if (id_filter >= 0 && aggregate->submission_id != (uint64_t)id_filter) {
      continue;
    }
    const iree_profile_projection_submission_timing_t timing =
        iree_profile_projection_calculate_submission_timing(
            context, aggregate->physical_device_ordinal,
            aggregate->earliest_start_tick, aggregate->latest_end_tick,
            aggregate->valid_count, aggregate->total_ticks);
    fprintf(file,
            "{\"type\":\"queue_submission\",\"submission_id\":%" PRIu64
            ",\"physical_device_ordinal\":%u,\"queue_ordinal\":%u"
            ",\"stream_id\":%" PRIu64 ",\"dispatches\":%" PRIu64
            ",\"valid\":%" PRIu64 ",\"invalid\":%" PRIu64
            ",\"span_ticks\":%.3f,\"total_dispatch_ticks\":%" PRIu64
            ",\"clock_fit_available\":%s"
            ",\"total_dispatch_time_ns_available\":%s"
            ",\"span_ns\":%.3f,\"total_dispatch_ns\":%" PRId64 "}\n",
            aggregate->submission_id, aggregate->physical_device_ordinal,
            aggregate->queue_ordinal, aggregate->stream_id,
            aggregate->dispatch_count, aggregate->valid_count,
            aggregate->invalid_count, timing.span_ticks, aggregate->total_ticks,
            timing.has_clock_fit ? "true" : "false",
            timing.has_total_dispatch_ns ? "true" : "false",
            timing.has_clock_fit ? timing.span_ticks * timing.ns_per_tick : 0.0,
            timing.has_total_dispatch_ns ? timing.total_dispatch_ns : 0);
  }
}

static void iree_profile_queue_print_jsonl_device_events(
    const iree_profile_dispatch_context_t* context, int64_t id_filter,
    FILE* file) {
  for (iree_host_size_t i = 0;
       i < context->queue_query.queue_device_event_count; ++i) {
    const iree_hal_profile_queue_device_event_t* event =
        &context->queue_query.queue_device_events[i].record;
    if (id_filter >= 0 && event->submission_id != (uint64_t)id_filter) {
      continue;
    }
    const bool is_valid = event->start_tick != 0 && event->end_tick != 0 &&
                          event->end_tick >= event->start_tick;
    const uint64_t duration_ticks =
        is_valid ? event->end_tick - event->start_tick : 0;
    iree_profile_model_clock_fit_t clock_fit;
    const bool has_clock_fit =
        iree_profile_projection_try_fit_driver_host_cpu_clock(
            context, event->physical_device_ordinal, &clock_fit);
    int64_t duration_ns = 0;
    const bool has_duration_ns = is_valid && has_clock_fit &&
                                 iree_profile_model_clock_fit_scale_ticks_to_ns(
                                     &clock_fit, duration_ticks, &duration_ns);
    fprintf(file,
            "{\"type\":\"queue_device_event\",\"event_id\":%" PRIu64
            ",\"op\":\"%s\",\"flags\":%u,\"submission_id\":%" PRIu64
            ",\"command_buffer_id\":%" PRIu64 ",\"allocation_id\":%" PRIu64
            ",\"physical_device_ordinal\":%u,\"queue_ordinal\":%u"
            ",\"stream_id\":%" PRIu64
            ",\"operation_count\":%u"
            ",\"payload_length\":%" PRIu64 ",\"start_tick\":%" PRIu64
            ",\"end_tick\":%" PRIu64 ",\"duration_ticks\":%" PRIu64
            ",\"valid\":%s"
            ",\"clock_fit_available\":%s,\"duration_ns\":%" PRId64 "}\n",
            event->event_id, iree_profile_queue_event_type_name(event->type),
            event->flags, event->submission_id, event->command_buffer_id,
            event->allocation_id, event->physical_device_ordinal,
            event->queue_ordinal, event->stream_id, event->operation_count,
            event->payload_length, event->start_tick, event->end_tick,
            duration_ticks, is_valid ? "true" : "false",
            has_clock_fit ? "true" : "false",
            has_duration_ns ? duration_ns : 0);
  }
}

static void iree_profile_queue_print_jsonl_host_execution_events(
    const iree_profile_dispatch_context_t* context, int64_t id_filter,
    FILE* file) {
  for (iree_host_size_t i = 0;
       i < context->queue_query.host_execution_event_count; ++i) {
    const iree_hal_profile_host_execution_event_t* event =
        &context->queue_query.host_execution_events[i].record;
    if (id_filter >= 0 && event->submission_id != (uint64_t)id_filter) {
      continue;
    }
    const bool is_valid = event->start_host_time_ns >= 0 &&
                          event->end_host_time_ns >= event->start_host_time_ns;
    const int64_t duration_ns =
        is_valid ? event->end_host_time_ns - event->start_host_time_ns : 0;
    fprintf(file,
            "{\"type\":\"host_execution_event\",\"event_id\":%" PRIu64
            ",\"op\":\"%s\",\"flags\":%u,\"status_code\":%u"
            ",\"submission_id\":%" PRIu64 ",\"command_buffer_id\":%" PRIu64
            ",\"command_index\":%u,\"executable_id\":%" PRIu64
            ",\"function_ordinal\":%u,\"allocation_id\":%" PRIu64
            ",\"physical_device_ordinal\":%u,\"queue_ordinal\":%u"
            ",\"stream_id\":%" PRIu64
            ",\"operation_count\":%u"
            ",\"payload_length\":%" PRIu64 ",\"tile_count\":%" PRIu64
            ",\"tile_duration_sum_ns\":%" PRId64
            ",\"host_time_domain\":\"iree_host_time_ns\""
            ",\"start_host_time_ns\":%" PRId64 ",\"end_host_time_ns\":%" PRId64
            ",\"duration_ns\":%" PRId64 ",\"valid\":%s}\n",
            event->event_id, iree_profile_queue_event_type_name(event->type),
            event->flags, event->status_code, event->submission_id,
            event->command_buffer_id, event->command_index,
            event->executable_id, event->function_ordinal, event->allocation_id,
            event->physical_device_ordinal, event->queue_ordinal,
            event->stream_id, event->operation_count, event->payload_length,
            event->tile_count, event->tile_duration_sum_ns,
            event->start_host_time_ns, event->end_host_time_ns, duration_ns,
            is_valid ? "true" : "false");
  }
}

static void iree_profile_queue_print_jsonl_events(
    const iree_profile_dispatch_context_t* context, int64_t id_filter,
    FILE* file) {
  for (iree_host_size_t i = 0; i < context->queue_query.queue_event_count;
       ++i) {
    const iree_hal_profile_queue_event_t* event =
        &context->queue_query.queue_events[i].record;
    if (id_filter >= 0 && event->submission_id != (uint64_t)id_filter) {
      continue;
    }
    fprintf(
        file,
        "{\"type\":\"queue_event\",\"event_id\":%" PRIu64
        ",\"op\":\"%s\",\"flags\":%u,\"dependency_strategy\":\"%s\""
        ",\"submission_id\":%" PRIu64 ",\"command_buffer_id\":%" PRIu64
        ",\"allocation_id\":%" PRIu64
        ",\"physical_device_ordinal\":%u,\"queue_ordinal\":%u"
        ",\"stream_id\":%" PRIu64 ",\"host_time_ns\":%" PRId64
        ",\"ready_host_time_ns\":%" PRId64
        ",\"host_time_domain\":\"iree_host_time_ns\""
        ",\"wait_count\":%u,\"signal_count\":%u,\"barrier_count\":%u"
        ",\"operation_count\":%u,\"payload_length\":%" PRIu64 "}\n",
        event->event_id, iree_profile_queue_event_type_name(event->type),
        event->flags,
        iree_profile_queue_dependency_strategy_name(event->dependency_strategy),
        event->submission_id, event->command_buffer_id, event->allocation_id,
        event->physical_device_ordinal, event->queue_ordinal, event->stream_id,
        event->host_time_ns, event->ready_host_time_ns, event->wait_count,
        event->signal_count, event->barrier_count, event->operation_count,
        event->payload_length);
  }
}

static iree_status_t iree_profile_queue_print_jsonl(
    const iree_profile_dispatch_context_t* context, iree_string_view_t filter,
    int64_t id_filter, FILE* file) {
  iree_profile_queue_print_jsonl_summary(context, filter, file);
  iree_profile_queue_print_jsonl_queues(context, file);
  iree_profile_queue_print_jsonl_submissions(context, id_filter, file);
  iree_profile_queue_print_jsonl_device_events(context, id_filter, file);
  iree_profile_queue_print_jsonl_host_execution_events(context, id_filter,
                                                       file);
  iree_profile_queue_print_jsonl_events(context, id_filter, file);
  return iree_ok_status();
}

typedef struct iree_profile_projection_parse_context_t {
  // Dispatch aggregation and shared model state for this projection.
  iree_profile_dispatch_context_t* dispatch_context;
  // Optional glob filter applied to projected keys.
  iree_string_view_t filter;
  // Projection selected by the command entry point.
  iree_profile_projection_mode_t projection_mode;
  // Optional entity identifier filter, or -1 when disabled.
  int64_t id_filter;
  // True when matched dispatches should update aggregate arrays.
  bool aggregate_events;
  // Optional callback receiving matched raw dispatch events.
  iree_profile_dispatch_event_callback_t event_callback;
} iree_profile_projection_parse_context_t;

static iree_status_t iree_profile_projection_metadata_record(
    void* user_data, const iree_hal_profile_file_record_t* record,
    iree_host_size_t record_index) {
  (void)record_index;
  iree_profile_projection_parse_context_t* context =
      (iree_profile_projection_parse_context_t*)user_data;
  return iree_profile_model_process_metadata_record(
      &context->dispatch_context->model, record);
}

static iree_status_t iree_profile_projection_event_record(
    void* user_data, const iree_hal_profile_file_record_t* record,
    iree_host_size_t record_index) {
  (void)record_index;
  iree_profile_projection_parse_context_t* context =
      (iree_profile_projection_parse_context_t*)user_data;
  return iree_profile_dispatch_process_events_record(
      context->dispatch_context, record, context->filter,
      context->projection_mode, context->id_filter, context->aggregate_events,
      context->event_callback);
}

static iree_status_t iree_profile_projection_print(
    const iree_profile_dispatch_context_t* context, iree_string_view_t filter,
    iree_profile_projection_mode_t projection_mode, int64_t id_filter,
    bool emit_events, bool is_text, FILE* file) {
  switch (projection_mode) {
    case IREE_PROFILE_PROJECTION_MODE_DISPATCH:
      if (is_text) {
        return iree_profile_dispatch_print_text(context, filter, file);
      }
      iree_profile_dispatch_print_jsonl_summary(context, filter, emit_events,
                                                file);
      return emit_events
                 ? iree_ok_status()
                 : iree_profile_dispatch_print_jsonl_aggregates(context, file);
    case IREE_PROFILE_PROJECTION_MODE_EXECUTABLE:
      if (emit_events) {
        iree_profile_dispatch_print_jsonl_summary(context, filter, emit_events,
                                                  file);
        return iree_ok_status();
      }
      {
        iree_profile_projection_index_t index;
        iree_profile_projection_index_initialize(context->host_allocator,
                                                 &index);
        iree_status_t status =
            iree_profile_projection_index_build_functions(context, &index);
        if (iree_status_is_ok(status)) {
          status = is_text ? iree_profile_executable_print_text(
                                 context, filter, id_filter, &index, file)
                           : iree_profile_executable_print_jsonl(
                                 context, filter, id_filter, &index, file);
        }
        iree_profile_projection_index_deinitialize(&index);
        return status;
      }
    case IREE_PROFILE_PROJECTION_MODE_COMMAND:
      if (emit_events) {
        iree_profile_dispatch_print_jsonl_summary(context, filter, emit_events,
                                                  file);
        return iree_ok_status();
      }
      {
        iree_profile_projection_index_t index;
        iree_profile_projection_index_initialize(context->host_allocator,
                                                 &index);
        iree_status_t status =
            iree_profile_projection_index_build_command_buffers(context,
                                                                &index);
        if (iree_status_is_ok(status)) {
          status = is_text ? iree_profile_command_print_text(
                                 context, filter, id_filter, &index, file)
                           : iree_profile_command_print_jsonl(
                                 context, filter, id_filter, &index, file);
        }
        iree_profile_projection_index_deinitialize(&index);
        return status;
      }
    case IREE_PROFILE_PROJECTION_MODE_QUEUE:
      if (emit_events) {
        iree_profile_dispatch_print_jsonl_summary(context, filter, emit_events,
                                                  file);
        return iree_ok_status();
      }
      if (!is_text) {
        return iree_profile_queue_print_jsonl(context, filter, id_filter, file);
      }
      {
        iree_profile_projection_index_t index;
        iree_profile_projection_index_initialize(context->host_allocator,
                                                 &index);
        iree_status_t status =
            iree_profile_projection_index_build_queues(context, &index);
        if (iree_status_is_ok(status)) {
          status = iree_profile_queue_print_text(context, filter, &index, file);
        }
        iree_profile_projection_index_deinitialize(&index);
        return status;
      }
    default:
      return iree_make_status(IREE_STATUS_INVALID_ARGUMENT,
                              "unsupported profile projection mode %d",
                              (int)projection_mode);
  }
}

iree_status_t iree_profile_projection_file(
    iree_string_view_t path, iree_string_view_t format,
    iree_string_view_t filter, iree_profile_projection_mode_t projection_mode,
    int64_t id_filter, bool emit_events, FILE* file,
    iree_allocator_t host_allocator) {
  bool is_text = iree_string_view_equal(format, IREE_SV("text"));
  bool is_jsonl = iree_string_view_equal(format, IREE_SV("jsonl"));
  if (!is_text && !is_jsonl) {
    return iree_make_status(IREE_STATUS_INVALID_ARGUMENT,
                            "unsupported profile output format '%.*s'",
                            (int)format.size, format.data);
  }
  if (emit_events && !is_jsonl) {
    return iree_make_status(IREE_STATUS_INVALID_ARGUMENT,
                            "--dispatch_events requires --format=jsonl");
  }

  iree_profile_file_t profile_file;
  IREE_RETURN_IF_ERROR(
      iree_profile_file_open(path, host_allocator, &profile_file));

  iree_profile_dispatch_context_t context;
  iree_profile_dispatch_context_initialize(host_allocator, &context);
  const iree_profile_dispatch_event_callback_t event_callback = {
      .fn = emit_events ? iree_profile_projection_emit_dispatch_event : NULL,
      .host_fn =
          emit_events ? iree_profile_projection_emit_host_dispatch_event : NULL,
      .user_data = file,
  };
  iree_profile_projection_parse_context_t parse_context = {
      .dispatch_context = &context,
      .filter = filter,
      .projection_mode = projection_mode,
      .id_filter = id_filter,
      .aggregate_events = !emit_events,
      .event_callback = event_callback,
  };
  iree_profile_file_record_callback_t record_callback = {
      .fn = iree_profile_projection_metadata_record,
      .user_data = &parse_context,
  };
  iree_status_t status =
      iree_profile_file_for_each_record(&profile_file, record_callback);
  if (iree_status_is_ok(status)) {
    record_callback.fn = iree_profile_projection_event_record;
    status = iree_profile_file_for_each_record(&profile_file, record_callback);
  }

  if (iree_status_is_ok(status)) {
    status =
        iree_profile_projection_print(&context, filter, projection_mode,
                                      id_filter, emit_events, is_text, file);
  }

  iree_profile_dispatch_context_deinitialize(&context);
  iree_profile_file_close(&profile_file);
  return status;
}
