// 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/model.h"

#include <string.h>

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

void iree_profile_model_initialize(iree_allocator_t host_allocator,
                                   iree_profile_model_t* out_model) {
  memset(out_model, 0, sizeof(*out_model));
  out_model->host_allocator = host_allocator;
}

void iree_profile_model_deinitialize(iree_profile_model_t* model) {
  iree_profile_index_deinitialize(&model->device_index, model->host_allocator);
  iree_profile_index_deinitialize(&model->metric_descriptor_index,
                                  model->host_allocator);
  iree_profile_index_deinitialize(&model->metric_source_index,
                                  model->host_allocator);
  iree_profile_index_deinitialize(&model->queue_index, model->host_allocator);
  iree_profile_index_deinitialize(&model->command_buffer_index,
                                  model->host_allocator);
  iree_profile_index_deinitialize(&model->function_index,
                                  model->host_allocator);
  iree_profile_index_deinitialize(&model->executable_index,
                                  model->host_allocator);
  iree_allocator_free(model->host_allocator, model->devices);
  iree_allocator_free(model->host_allocator, model->metric_descriptors);
  iree_allocator_free(model->host_allocator, model->metric_sources);
  iree_allocator_free(model->host_allocator, model->queues);
  iree_allocator_free(model->host_allocator, model->command_operations);
  iree_allocator_free(model->host_allocator, model->command_buffers);
  iree_allocator_free(model->host_allocator, model->functions);
  iree_allocator_free(model->host_allocator, model->executables);
  memset(model, 0, sizeof(*model));
}

typedef struct iree_profile_model_device_lookup_t {
  // Model owning the candidate rows.
  const iree_profile_model_t* model;
  // Session-local physical device ordinal.
  uint32_t physical_device_ordinal;
} iree_profile_model_device_lookup_t;

typedef struct iree_profile_model_queue_lookup_t {
  // Model owning the candidate rows.
  const iree_profile_model_t* model;
  // Session-local physical device ordinal.
  uint32_t physical_device_ordinal;
  // Session-local queue ordinal.
  uint32_t queue_ordinal;
  // Producer-defined queue stream identifier.
  uint64_t stream_id;
} iree_profile_model_queue_lookup_t;

typedef struct iree_profile_model_executable_lookup_t {
  // Model owning the candidate rows.
  const iree_profile_model_t* model;
  // Producer-local executable identifier.
  uint64_t executable_id;
} iree_profile_model_executable_lookup_t;

typedef struct iree_profile_model_function_lookup_t {
  // Model owning the candidate rows.
  const iree_profile_model_t* model;
  // Producer-local executable identifier.
  uint64_t executable_id;
  // Function ordinal within |executable_id|.
  uint32_t function_ordinal;
} iree_profile_model_function_lookup_t;

typedef struct iree_profile_model_command_buffer_lookup_t {
  // Model owning the candidate rows.
  const iree_profile_model_t* model;
  // Producer-local command-buffer identifier.
  uint64_t command_buffer_id;
} iree_profile_model_command_buffer_lookup_t;

typedef struct iree_profile_model_metric_source_lookup_t {
  // Model owning the candidate rows.
  const iree_profile_model_t* model;
  // Producer-defined metric source identifier.
  uint64_t source_id;
} iree_profile_model_metric_source_lookup_t;

typedef struct iree_profile_model_metric_descriptor_lookup_t {
  // Model owning the candidate rows.
  const iree_profile_model_t* model;
  // Producer-defined metric source identifier.
  uint64_t source_id;
  // Source-scoped metric identifier.
  uint64_t metric_id;
} iree_profile_model_metric_descriptor_lookup_t;

static uint64_t iree_profile_model_device_hash(
    uint32_t physical_device_ordinal) {
  return iree_profile_index_mix_u64(physical_device_ordinal);
}

static uint64_t iree_profile_model_queue_hash(uint32_t physical_device_ordinal,
                                              uint32_t queue_ordinal,
                                              uint64_t stream_id) {
  uint64_t hash = iree_profile_model_device_hash(physical_device_ordinal);
  hash = iree_profile_index_combine_u64(hash, queue_ordinal);
  return iree_profile_index_combine_u64(hash, stream_id);
}

static uint64_t iree_profile_model_executable_hash(uint64_t executable_id) {
  return iree_profile_index_mix_u64(executable_id);
}

static uint64_t iree_profile_model_function_hash(uint64_t executable_id,
                                                 uint32_t function_ordinal) {
  uint64_t hash = iree_profile_model_executable_hash(executable_id);
  return iree_profile_index_combine_u64(hash, function_ordinal);
}

static uint64_t iree_profile_model_command_buffer_hash(
    uint64_t command_buffer_id) {
  return iree_profile_index_mix_u64(command_buffer_id);
}

static uint64_t iree_profile_model_metric_source_hash(uint64_t source_id) {
  return iree_profile_index_mix_u64(source_id);
}

static uint64_t iree_profile_model_metric_descriptor_hash(uint64_t source_id,
                                                          uint64_t metric_id) {
  uint64_t hash = iree_profile_model_metric_source_hash(source_id);
  return iree_profile_index_combine_u64(hash, metric_id);
}

static bool iree_profile_model_device_matches(const void* user_data,
                                              iree_host_size_t value) {
  const iree_profile_model_device_lookup_t* lookup =
      (const iree_profile_model_device_lookup_t*)user_data;
  return lookup->model->devices[value].physical_device_ordinal ==
         lookup->physical_device_ordinal;
}

static bool iree_profile_model_queue_matches(const void* user_data,
                                             iree_host_size_t value) {
  const iree_profile_model_queue_lookup_t* lookup =
      (const iree_profile_model_queue_lookup_t*)user_data;
  const iree_hal_profile_queue_record_t* record =
      &lookup->model->queues[value].record;
  return record->physical_device_ordinal == lookup->physical_device_ordinal &&
         record->queue_ordinal == lookup->queue_ordinal &&
         record->stream_id == lookup->stream_id;
}

static bool iree_profile_model_executable_matches(const void* user_data,
                                                  iree_host_size_t value) {
  const iree_profile_model_executable_lookup_t* lookup =
      (const iree_profile_model_executable_lookup_t*)user_data;
  return lookup->model->executables[value].record.executable_id ==
         lookup->executable_id;
}

static bool iree_profile_model_function_matches(const void* user_data,
                                                iree_host_size_t value) {
  const iree_profile_model_function_lookup_t* lookup =
      (const iree_profile_model_function_lookup_t*)user_data;
  const iree_profile_model_function_t* function_info =
      &lookup->model->functions[value];
  return function_info->executable_id == lookup->executable_id &&
         function_info->function_ordinal == lookup->function_ordinal;
}

static bool iree_profile_model_command_buffer_matches(const void* user_data,
                                                      iree_host_size_t value) {
  const iree_profile_model_command_buffer_lookup_t* lookup =
      (const iree_profile_model_command_buffer_lookup_t*)user_data;
  return lookup->model->command_buffers[value].record.command_buffer_id ==
         lookup->command_buffer_id;
}

static bool iree_profile_model_metric_source_matches(const void* user_data,
                                                     iree_host_size_t value) {
  const iree_profile_model_metric_source_lookup_t* lookup =
      (const iree_profile_model_metric_source_lookup_t*)user_data;
  return lookup->model->metric_sources[value].record.source_id ==
         lookup->source_id;
}

static bool iree_profile_model_metric_descriptor_matches(
    const void* user_data, iree_host_size_t value) {
  const iree_profile_model_metric_descriptor_lookup_t* lookup =
      (const iree_profile_model_metric_descriptor_lookup_t*)user_data;
  const iree_hal_profile_device_metric_descriptor_record_t* record =
      &lookup->model->metric_descriptors[value].record;
  return record->source_id == lookup->source_id &&
         record->metric_id == lookup->metric_id;
}

iree_status_t iree_profile_model_ensure_device(
    iree_profile_model_t* model, uint32_t physical_device_ordinal,
    iree_profile_model_device_t** out_device) {
  *out_device = NULL;

  const iree_profile_model_device_lookup_t lookup = {
      .model = model,
      .physical_device_ordinal = physical_device_ordinal,
  };
  const uint64_t hash = iree_profile_model_device_hash(physical_device_ordinal);
  iree_host_size_t existing_index = 0;
  if (iree_profile_index_find(&model->device_index, hash,
                              iree_profile_model_device_matches, &lookup,
                              &existing_index)) {
    *out_device = &model->devices[existing_index];
    return iree_ok_status();
  }

  if (model->device_count + 1 > model->device_capacity) {
    IREE_RETURN_IF_ERROR(iree_allocator_grow_array(
        model->host_allocator,
        iree_max((iree_host_size_t)4, model->device_count + 1),
        sizeof(model->devices[0]), &model->device_capacity,
        (void**)&model->devices));
  }

  const iree_host_size_t device_index = model->device_count;
  IREE_RETURN_IF_ERROR(iree_profile_index_insert(
      &model->device_index, model->host_allocator, hash, device_index));
  ++model->device_count;

  iree_profile_model_device_t* device = &model->devices[device_index];
  memset(device, 0, sizeof(*device));
  device->physical_device_ordinal = physical_device_ordinal;
  *out_device = device;
  return iree_ok_status();
}

static void iree_profile_model_record_clock_sample(
    iree_profile_model_device_t* device,
    const iree_hal_profile_clock_correlation_record_t* record) {
  if (iree_any_bit_set(
          record->flags,
          IREE_HAL_PROFILE_CLOCK_CORRELATION_FLAG_DEVICE_TICK_UNALIGNED)) {
    ++device->invalid_clock_alignment_sample_count;
  }
  if (device->clock_sample_count == 0) {
    device->first_clock_sample = *record;
  }
  device->last_clock_sample = *record;
  ++device->clock_sample_count;
}

const iree_profile_model_device_t* iree_profile_model_find_device(
    const iree_profile_model_t* model, uint32_t physical_device_ordinal) {
  if (!model) return NULL;
  const iree_profile_model_device_lookup_t lookup = {
      .model = model,
      .physical_device_ordinal = physical_device_ordinal,
  };
  iree_host_size_t index = 0;
  if (iree_profile_index_find(
          &model->device_index,
          iree_profile_model_device_hash(physical_device_ordinal),
          iree_profile_model_device_matches, &lookup, &index)) {
    return &model->devices[index];
  }
  return NULL;
}

const iree_profile_model_queue_t* iree_profile_model_find_queue(
    const iree_profile_model_t* model, uint32_t physical_device_ordinal,
    uint32_t queue_ordinal, uint64_t stream_id) {
  if (!model) return NULL;
  const iree_profile_model_queue_lookup_t lookup = {
      .model = model,
      .physical_device_ordinal = physical_device_ordinal,
      .queue_ordinal = queue_ordinal,
      .stream_id = stream_id,
  };
  iree_host_size_t index = 0;
  if (iree_profile_index_find(
          &model->queue_index,
          iree_profile_model_queue_hash(physical_device_ordinal, queue_ordinal,
                                        stream_id),
          iree_profile_model_queue_matches, &lookup, &index)) {
    return &model->queues[index];
  }
  return NULL;
}

const iree_profile_model_executable_t* iree_profile_model_find_executable(
    const iree_profile_model_t* model, uint64_t executable_id) {
  if (!model) return NULL;
  const iree_profile_model_executable_lookup_t lookup = {
      .model = model,
      .executable_id = executable_id,
  };
  iree_host_size_t index = 0;
  if (iree_profile_index_find(&model->executable_index,
                              iree_profile_model_executable_hash(executable_id),
                              iree_profile_model_executable_matches, &lookup,
                              &index)) {
    return &model->executables[index];
  }
  return NULL;
}

static iree_profile_model_executable_t*
iree_profile_model_find_executable_mutable(iree_profile_model_t* model,
                                           uint64_t executable_id) {
  if (!model) return NULL;
  const iree_profile_model_executable_lookup_t lookup = {
      .model = model,
      .executable_id = executable_id,
  };
  iree_host_size_t index = 0;
  if (iree_profile_index_find(&model->executable_index,
                              iree_profile_model_executable_hash(executable_id),
                              iree_profile_model_executable_matches, &lookup,
                              &index)) {
    return &model->executables[index];
  }
  return NULL;
}

const iree_profile_model_function_t* iree_profile_model_find_function(
    const iree_profile_model_t* model, uint64_t executable_id,
    uint32_t function_ordinal) {
  if (!model) return NULL;
  const iree_profile_model_function_lookup_t lookup = {
      .model = model,
      .executable_id = executable_id,
      .function_ordinal = function_ordinal,
  };
  iree_host_size_t index = 0;
  if (iree_profile_index_find(
          &model->function_index,
          iree_profile_model_function_hash(executable_id, function_ordinal),
          iree_profile_model_function_matches, &lookup, &index)) {
    return &model->functions[index];
  }
  return NULL;
}

const iree_profile_model_command_buffer_t*
iree_profile_model_find_command_buffer(const iree_profile_model_t* model,
                                       uint64_t command_buffer_id) {
  if (!model) return NULL;
  const iree_profile_model_command_buffer_lookup_t lookup = {
      .model = model,
      .command_buffer_id = command_buffer_id,
  };
  iree_host_size_t index = 0;
  if (iree_profile_index_find(
          &model->command_buffer_index,
          iree_profile_model_command_buffer_hash(command_buffer_id),
          iree_profile_model_command_buffer_matches, &lookup, &index)) {
    return &model->command_buffers[index];
  }
  return NULL;
}

const iree_profile_model_metric_source_t* iree_profile_model_find_metric_source(
    const iree_profile_model_t* model, uint64_t source_id) {
  if (!model) return NULL;
  const iree_profile_model_metric_source_lookup_t lookup = {
      .model = model,
      .source_id = source_id,
  };
  iree_host_size_t index = 0;
  if (iree_profile_index_find(&model->metric_source_index,
                              iree_profile_model_metric_source_hash(source_id),
                              iree_profile_model_metric_source_matches, &lookup,
                              &index)) {
    return &model->metric_sources[index];
  }
  return NULL;
}

const iree_profile_model_metric_descriptor_t*
iree_profile_model_find_metric_descriptor(const iree_profile_model_t* model,
                                          uint64_t source_id,
                                          uint64_t metric_id) {
  if (!model) return NULL;
  const iree_profile_model_metric_descriptor_lookup_t lookup = {
      .model = model,
      .source_id = source_id,
      .metric_id = metric_id,
  };
  iree_host_size_t index = 0;
  if (iree_profile_index_find(
          &model->metric_descriptor_index,
          iree_profile_model_metric_descriptor_hash(source_id, metric_id),
          iree_profile_model_metric_descriptor_matches, &lookup, &index)) {
    return &model->metric_descriptors[index];
  }
  return NULL;
}

iree_status_t iree_profile_model_resolve_metric_descriptor(
    const iree_profile_model_t* model, uint64_t source_id, uint64_t metric_id,
    const iree_profile_model_metric_descriptor_t** out_descriptor) {
  *out_descriptor =
      iree_profile_model_find_metric_descriptor(model, source_id, metric_id);
  if (!*out_descriptor) {
    return iree_make_status(
        IREE_STATUS_DATA_LOSS,
        "device metric sample references missing descriptor source=%" PRIu64
        " metric=%" PRIu64,
        source_id, metric_id);
  }
  return iree_ok_status();
}

static iree_profile_model_command_buffer_t*
iree_profile_model_find_command_buffer_mutable(iree_profile_model_t* model,
                                               uint64_t command_buffer_id) {
  if (!model) return NULL;
  const iree_profile_model_command_buffer_lookup_t lookup = {
      .model = model,
      .command_buffer_id = command_buffer_id,
  };
  iree_host_size_t index = 0;
  if (iree_profile_index_find(
          &model->command_buffer_index,
          iree_profile_model_command_buffer_hash(command_buffer_id),
          iree_profile_model_command_buffer_matches, &lookup, &index)) {
    return &model->command_buffers[index];
  }
  return NULL;
}

static bool iree_profile_model_host_time_midpoint(
    const iree_hal_profile_clock_correlation_record_t* sample,
    int64_t* out_time_ns) {
  if (!iree_all_bits_set(
          sample->flags,
          IREE_HAL_PROFILE_CLOCK_CORRELATION_FLAG_HOST_TIME_BRACKET) ||
      sample->host_time_begin_ns < 0 ||
      sample->host_time_end_ns < sample->host_time_begin_ns) {
    return false;
  }
  *out_time_ns = sample->host_time_begin_ns +
                 (sample->host_time_end_ns - sample->host_time_begin_ns) / 2;
  return true;
}

static bool iree_profile_model_sample_time_ns(
    const iree_hal_profile_clock_correlation_record_t* sample,
    iree_profile_model_clock_time_domain_t time_domain, int64_t* out_time_ns) {
  switch (time_domain) {
    case IREE_PROFILE_MODEL_CLOCK_TIME_DOMAIN_HOST_CPU_TIMESTAMP_NS:
      if (!iree_all_bits_set(
              sample->flags,
              IREE_HAL_PROFILE_CLOCK_CORRELATION_FLAG_HOST_CPU_TIMESTAMP) ||
          sample->host_cpu_timestamp_ns > INT64_MAX) {
        return false;
      }
      *out_time_ns = (int64_t)sample->host_cpu_timestamp_ns;
      return true;
    case IREE_PROFILE_MODEL_CLOCK_TIME_DOMAIN_IREE_HOST_TIME_NS:
      return iree_profile_model_host_time_midpoint(sample, out_time_ns);
    default:
      return false;
  }
}

bool iree_profile_model_device_try_fit_clock_exact(
    const iree_profile_model_device_t* device,
    iree_profile_model_clock_time_domain_t time_domain,
    iree_profile_model_clock_fit_t* out_fit) {
  memset(out_fit, 0, sizeof(*out_fit));
  if (!device || device->clock_sample_count < 2) return false;
  if (device->invalid_clock_alignment_sample_count != 0) return false;

  const iree_hal_profile_clock_correlation_record_t* first =
      &device->first_clock_sample;
  const iree_hal_profile_clock_correlation_record_t* last =
      &device->last_clock_sample;
  if (!iree_all_bits_set(first->flags,
                         IREE_HAL_PROFILE_CLOCK_CORRELATION_FLAG_DEVICE_TICK) ||
      !iree_all_bits_set(last->flags,
                         IREE_HAL_PROFILE_CLOCK_CORRELATION_FLAG_DEVICE_TICK) ||
      iree_any_bit_set(
          first->flags | last->flags,
          IREE_HAL_PROFILE_CLOCK_CORRELATION_FLAG_DEVICE_TICK_UNALIGNED)) {
    return false;
  }
  int64_t first_time_ns = 0;
  int64_t last_time_ns = 0;
  if (!iree_profile_model_sample_time_ns(first, time_domain, &first_time_ns) ||
      !iree_profile_model_sample_time_ns(last, time_domain, &last_time_ns)) {
    return false;
  }
  if (last->device_tick <= first->device_tick ||
      last_time_ns <= first_time_ns) {
    return false;
  }

  out_fit->time_domain = time_domain;
  out_fit->first_sample_id = first->sample_id;
  out_fit->last_sample_id = last->sample_id;
  out_fit->first_device_tick = first->device_tick;
  out_fit->last_device_tick = last->device_tick;
  out_fit->first_time_ns = first_time_ns;
  out_fit->last_time_ns = last_time_ns;
  out_fit->device_tick_span = last->device_tick - first->device_tick;
  out_fit->time_span_ns = (uint64_t)(last_time_ns - first_time_ns);
  return true;
}

static bool iree_profile_model_round_mul_div_u64(uint64_t value,
                                                 uint64_t numerator,
                                                 uint64_t denominator,
                                                 uint64_t* out_result) {
  *out_result = 0;
  if (denominator == 0) return false;
  if (value == 0 || numerator == 0) return true;

#if defined(__SIZEOF_INT128__)
  __uint128_t product = (__uint128_t)value * (__uint128_t)numerator;
  product += denominator / 2;
  __uint128_t quotient = product / denominator;
  if (quotient > UINT64_MAX) return false;
  *out_result = (uint64_t)quotient;
  return true;
#else
  const uint64_t whole = value / denominator;
  const uint64_t remainder = value % denominator;
  if (whole > UINT64_MAX / numerator) return false;
  uint64_t scaled = whole * numerator;
  if (remainder != 0) {
    if (remainder > UINT64_MAX / numerator) return false;
    uint64_t fractional_product = remainder * numerator;
    if (fractional_product > UINT64_MAX - denominator / 2) return false;
    uint64_t fractional = (fractional_product + denominator / 2) / denominator;
    if (scaled > UINT64_MAX - fractional) return false;
    scaled += fractional;
  }
  *out_result = scaled;
  return true;
#endif  // defined(__SIZEOF_INT128__)
}

bool iree_profile_model_clock_fit_scale_ticks_to_ns(
    const iree_profile_model_clock_fit_t* fit, uint64_t device_tick_count,
    int64_t* out_duration_ns) {
  *out_duration_ns = 0;
  if (!fit || fit->device_tick_span == 0) return false;
  uint64_t duration_ns = 0;
  if (!iree_profile_model_round_mul_div_u64(
          device_tick_count, fit->time_span_ns, fit->device_tick_span,
          &duration_ns) ||
      duration_ns > INT64_MAX) {
    return false;
  }
  *out_duration_ns = (int64_t)duration_ns;
  return true;
}

bool iree_profile_model_clock_fit_map_tick(
    const iree_profile_model_clock_fit_t* fit, uint64_t device_tick,
    int64_t* out_time_ns) {
  *out_time_ns = 0;
  if (!fit || fit->device_tick_span == 0) return false;

  const bool after_base = device_tick >= fit->first_device_tick;
  const uint64_t tick_delta = after_base ? device_tick - fit->first_device_tick
                                         : fit->first_device_tick - device_tick;
  int64_t time_delta_ns = 0;
  if (!iree_profile_model_clock_fit_scale_ticks_to_ns(fit, tick_delta,
                                                      &time_delta_ns)) {
    return false;
  }

  if (after_base) {
    if (fit->first_time_ns > INT64_MAX - time_delta_ns) return false;
    *out_time_ns = fit->first_time_ns + time_delta_ns;
  } else {
    if (fit->first_time_ns < INT64_MIN + time_delta_ns) return false;
    *out_time_ns = fit->first_time_ns - time_delta_ns;
  }
  return true;
}

double iree_profile_model_clock_fit_ns_per_tick(
    const iree_profile_model_clock_fit_t* fit) {
  if (!fit || fit->device_tick_span == 0) return 0.0;
  return (double)fit->time_span_ns / (double)fit->device_tick_span;
}

double iree_profile_model_clock_fit_tick_frequency_hz(
    const iree_profile_model_clock_fit_t* fit) {
  const double ns_per_tick = iree_profile_model_clock_fit_ns_per_tick(fit);
  return ns_per_tick > 0.0 ? 1000000000.0 / ns_per_tick : 0.0;
}

bool iree_profile_model_device_try_fit_clock(
    const iree_profile_model_device_t* device, double* out_ns_per_tick,
    double* out_tick_frequency_hz) {
  *out_ns_per_tick = 0.0;
  *out_tick_frequency_hz = 0.0;
  iree_profile_model_clock_fit_t fit;
  if (!iree_profile_model_device_try_fit_clock_exact(
          device, IREE_PROFILE_MODEL_CLOCK_TIME_DOMAIN_HOST_CPU_TIMESTAMP_NS,
          &fit)) {
    return false;
  }
  *out_ns_per_tick = iree_profile_model_clock_fit_ns_per_tick(&fit);
  *out_tick_frequency_hz = iree_profile_model_clock_fit_tick_frequency_hz(&fit);
  return true;
}

static iree_string_view_t iree_profile_model_format_numeric_key(
    uint32_t physical_device_ordinal, uint64_t executable_id,
    uint32_t function_ordinal, char* buffer, iree_host_size_t buffer_capacity) {
  if (buffer_capacity == 0) return iree_string_view_empty();
  int result = 0;
  if (physical_device_ordinal == UINT32_MAX) {
    result = snprintf(buffer, buffer_capacity, "executable%" PRIu64 "#%u",
                      executable_id, function_ordinal);
  } else {
    result =
        snprintf(buffer, buffer_capacity, "device%u/executable%" PRIu64 "#%u",
                 physical_device_ordinal, executable_id, function_ordinal);
  }
  if (result < 0) return iree_string_view_empty();
  iree_host_size_t length = (iree_host_size_t)result;
  if (length >= buffer_capacity) length = buffer_capacity - 1;
  return iree_make_string_view(buffer, length);
}

iree_string_view_t iree_profile_model_format_function_key(
    const iree_profile_model_function_t* function_info,
    uint32_t physical_device_ordinal, char* numeric_buffer,
    iree_host_size_t numeric_buffer_capacity) {
  if (!iree_string_view_is_empty(function_info->name)) {
    return function_info->name;
  }
  return iree_profile_model_format_numeric_key(
      physical_device_ordinal, function_info->executable_id,
      function_info->function_ordinal, numeric_buffer, numeric_buffer_capacity);
}

iree_status_t iree_profile_model_resolve_dispatch_key(
    const iree_profile_model_t* model, uint32_t physical_device_ordinal,
    uint64_t executable_id, uint32_t function_ordinal, char* numeric_buffer,
    iree_host_size_t numeric_buffer_capacity, iree_string_view_t* out_key) {
  *out_key = iree_string_view_empty();
  const iree_profile_model_executable_t* executable_info =
      iree_profile_model_find_executable(model, executable_id);
  if (!executable_info) {
    return iree_make_status(
        IREE_STATUS_DATA_LOSS,
        "dispatch event references missing executable metadata "
        "device=%u executable=%" PRIu64 " function=%u",
        physical_device_ordinal, executable_id, function_ordinal);
  }
  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,
        "dispatch event references missing executable function metadata "
        "device=%u executable=%" PRIu64 " function=%u",
        physical_device_ordinal, executable_id, function_ordinal);
  }
  *out_key = iree_profile_model_format_function_key(
      function_info, physical_device_ordinal, numeric_buffer,
      numeric_buffer_capacity);
  return iree_ok_status();
}

static iree_status_t iree_profile_model_append_function(
    iree_profile_model_t* model,
    const iree_profile_model_function_t* function_info) {
  iree_profile_model_executable_t* executable =
      iree_profile_model_find_executable_mutable(model,
                                                 function_info->executable_id);
  if (!executable) {
    return iree_make_status(
        IREE_STATUS_DATA_LOSS,
        "executable function references missing executable metadata "
        "executable=%" PRIu64 " function=%u",
        function_info->executable_id, function_info->function_ordinal);
  }

  const iree_profile_model_function_lookup_t lookup = {
      .model = model,
      .executable_id = function_info->executable_id,
      .function_ordinal = function_info->function_ordinal,
  };
  const uint64_t hash = iree_profile_model_function_hash(
      function_info->executable_id, function_info->function_ordinal);
  iree_host_size_t existing_index = 0;
  if (iree_profile_index_find(&model->function_index, hash,
                              iree_profile_model_function_matches, &lookup,
                              &existing_index)) {
    return iree_ok_status();
  }

  if (model->function_count + 1 > model->function_capacity) {
    IREE_RETURN_IF_ERROR(iree_allocator_grow_array(
        model->host_allocator,
        iree_max((iree_host_size_t)16, model->function_count + 1),
        sizeof(model->functions[0]), &model->function_capacity,
        (void**)&model->functions));
  }

  const iree_host_size_t function_index = model->function_count;
  IREE_RETURN_IF_ERROR(iree_profile_index_insert(
      &model->function_index, model->host_allocator, hash, function_index));
  model->functions[function_index] = *function_info;
  model->functions[function_index].next_function_index = IREE_HOST_SIZE_MAX;
  if (executable->first_function_index == IREE_HOST_SIZE_MAX) {
    executable->first_function_index = function_index;
  } else {
    model->functions[executable->last_function_index].next_function_index =
        function_index;
  }
  executable->last_function_index = function_index;
  ++executable->function_row_count;
  ++model->function_count;
  return iree_ok_status();
}

static iree_status_t iree_profile_model_append_executable(
    iree_profile_model_t* model,
    const iree_hal_profile_executable_record_t* record) {
  const iree_profile_model_executable_lookup_t lookup = {
      .model = model,
      .executable_id = record->executable_id,
  };
  const uint64_t hash =
      iree_profile_model_executable_hash(record->executable_id);
  iree_host_size_t existing_index = 0;
  if (iree_profile_index_find(&model->executable_index, hash,
                              iree_profile_model_executable_matches, &lookup,
                              &existing_index)) {
    return iree_ok_status();
  }

  if (model->executable_count + 1 > model->executable_capacity) {
    IREE_RETURN_IF_ERROR(iree_allocator_grow_array(
        model->host_allocator,
        iree_max((iree_host_size_t)8, model->executable_count + 1),
        sizeof(model->executables[0]), &model->executable_capacity,
        (void**)&model->executables));
  }
  const iree_host_size_t executable_index = model->executable_count;
  IREE_RETURN_IF_ERROR(iree_profile_index_insert(
      &model->executable_index, model->host_allocator, hash, executable_index));
  iree_profile_model_executable_t* executable_info =
      &model->executables[executable_index];
  executable_info->record = *record;
  executable_info->first_function_index = IREE_HOST_SIZE_MAX;
  executable_info->last_function_index = IREE_HOST_SIZE_MAX;
  executable_info->function_row_count = 0;
  ++model->executable_count;
  return iree_ok_status();
}

static iree_status_t iree_profile_model_append_command_buffer(
    iree_profile_model_t* model,
    const iree_hal_profile_command_buffer_record_t* record) {
  const iree_profile_model_command_buffer_lookup_t lookup = {
      .model = model,
      .command_buffer_id = record->command_buffer_id,
  };
  const uint64_t hash =
      iree_profile_model_command_buffer_hash(record->command_buffer_id);
  iree_host_size_t existing_index = 0;
  if (iree_profile_index_find(&model->command_buffer_index, hash,
                              iree_profile_model_command_buffer_matches,
                              &lookup, &existing_index)) {
    return iree_ok_status();
  }

  if (model->command_buffer_count + 1 > model->command_buffer_capacity) {
    IREE_RETURN_IF_ERROR(iree_allocator_grow_array(
        model->host_allocator,
        iree_max((iree_host_size_t)16, model->command_buffer_count + 1),
        sizeof(model->command_buffers[0]), &model->command_buffer_capacity,
        (void**)&model->command_buffers));
  }
  const iree_host_size_t command_buffer_index = model->command_buffer_count;
  IREE_RETURN_IF_ERROR(iree_profile_index_insert(&model->command_buffer_index,
                                                 model->host_allocator, hash,
                                                 command_buffer_index));
  iree_profile_model_command_buffer_t* command_buffer_info =
      &model->command_buffers[command_buffer_index];
  command_buffer_info->record = *record;
  command_buffer_info->first_operation_index = IREE_HOST_SIZE_MAX;
  command_buffer_info->last_operation_index = IREE_HOST_SIZE_MAX;
  command_buffer_info->operation_count = 0;
  ++model->command_buffer_count;
  return iree_ok_status();
}

static iree_status_t iree_profile_model_append_command_operation(
    iree_profile_model_t* model,
    const iree_hal_profile_command_operation_record_t* record) {
  iree_profile_model_command_buffer_t* command_buffer =
      iree_profile_model_find_command_buffer_mutable(model,
                                                     record->command_buffer_id);
  if (!command_buffer) {
    return iree_make_status(
        IREE_STATUS_DATA_LOSS,
        "command operation references missing command-buffer metadata "
        "command_buffer=%" PRIu64 " command_index=%u",
        record->command_buffer_id, record->command_index);
  }

  if (model->command_operation_count + 1 > model->command_operation_capacity) {
    IREE_RETURN_IF_ERROR(iree_allocator_grow_array(
        model->host_allocator,
        iree_max((iree_host_size_t)64, model->command_operation_count + 1),
        sizeof(model->command_operations[0]),
        &model->command_operation_capacity,
        (void**)&model->command_operations));
  }

  const iree_host_size_t operation_index = model->command_operation_count;
  iree_profile_model_command_operation_t* operation_info =
      &model->command_operations[operation_index];
  operation_info->record = *record;
  operation_info->next_operation_index = IREE_HOST_SIZE_MAX;
  if (command_buffer->first_operation_index == IREE_HOST_SIZE_MAX) {
    command_buffer->first_operation_index = operation_index;
  } else {
    model->command_operations[command_buffer->last_operation_index]
        .next_operation_index = operation_index;
  }
  command_buffer->last_operation_index = operation_index;
  ++command_buffer->operation_count;
  ++model->command_operation_count;
  return iree_ok_status();
}

static iree_status_t iree_profile_model_validate_command_operation(
    const iree_hal_profile_command_operation_record_t* record) {
  const bool has_block_structure =
      iree_hal_profile_command_operation_has_block_structure(record);
  if (has_block_structure) {
    if (record->block_ordinal == UINT32_MAX ||
        record->block_command_ordinal == UINT32_MAX) {
      return iree_make_status(
          IREE_STATUS_DATA_LOSS,
          "command operation declares block structure without block "
          "coordinates command_buffer=%" PRIu64 " command_index=%u",
          record->command_buffer_id, record->command_index);
    }
  } else if (record->block_ordinal != UINT32_MAX ||
             record->block_command_ordinal != UINT32_MAX ||
             record->target_block_ordinal != UINT32_MAX ||
             record->alternate_block_ordinal != UINT32_MAX) {
    return iree_make_status(
        IREE_STATUS_DATA_LOSS,
        "command operation carries block coordinates without the block "
        "structure flag command_buffer=%" PRIu64 " command_index=%u",
        record->command_buffer_id, record->command_index);
  }
  return iree_ok_status();
}

static iree_status_t iree_profile_model_append_queue(
    iree_profile_model_t* model,
    const iree_hal_profile_queue_record_t* record) {
  const iree_profile_model_queue_lookup_t lookup = {
      .model = model,
      .physical_device_ordinal = record->physical_device_ordinal,
      .queue_ordinal = record->queue_ordinal,
      .stream_id = record->stream_id,
  };
  const uint64_t hash =
      iree_profile_model_queue_hash(record->physical_device_ordinal,
                                    record->queue_ordinal, record->stream_id);
  iree_host_size_t existing_index = 0;
  if (iree_profile_index_find(&model->queue_index, hash,
                              iree_profile_model_queue_matches, &lookup,
                              &existing_index)) {
    return iree_ok_status();
  }

  if (model->queue_count + 1 > model->queue_capacity) {
    IREE_RETURN_IF_ERROR(iree_allocator_grow_array(
        model->host_allocator,
        iree_max((iree_host_size_t)4, model->queue_count + 1),
        sizeof(model->queues[0]), &model->queue_capacity,
        (void**)&model->queues));
  }
  const iree_host_size_t queue_index = model->queue_count;
  IREE_RETURN_IF_ERROR(iree_profile_index_insert(
      &model->queue_index, model->host_allocator, hash, queue_index));
  iree_profile_model_queue_t* queue_info = &model->queues[queue_index];
  queue_info->record = *record;
  ++model->queue_count;
  return iree_ok_status();
}

static iree_status_t iree_profile_model_append_metric_source(
    iree_profile_model_t* model,
    const iree_profile_model_metric_source_t* source_info) {
  const iree_profile_model_metric_source_lookup_t lookup = {
      .model = model,
      .source_id = source_info->record.source_id,
  };
  const uint64_t hash =
      iree_profile_model_metric_source_hash(source_info->record.source_id);
  iree_host_size_t existing_index = 0;
  if (iree_profile_index_find(&model->metric_source_index, hash,
                              iree_profile_model_metric_source_matches, &lookup,
                              &existing_index)) {
    return iree_ok_status();
  }

  if (model->metric_source_count + 1 > model->metric_source_capacity) {
    IREE_RETURN_IF_ERROR(iree_allocator_grow_array(
        model->host_allocator,
        iree_max((iree_host_size_t)4, model->metric_source_count + 1),
        sizeof(model->metric_sources[0]), &model->metric_source_capacity,
        (void**)&model->metric_sources));
  }
  const iree_host_size_t source_index = model->metric_source_count;
  IREE_RETURN_IF_ERROR(iree_profile_index_insert(
      &model->metric_source_index, model->host_allocator, hash, source_index));
  model->metric_sources[source_index] = *source_info;
  ++model->metric_source_count;
  return iree_ok_status();
}

static iree_status_t iree_profile_model_append_metric_descriptor(
    iree_profile_model_t* model,
    const iree_profile_model_metric_descriptor_t* descriptor_info) {
  const iree_profile_model_metric_descriptor_lookup_t lookup = {
      .model = model,
      .source_id = descriptor_info->record.source_id,
      .metric_id = descriptor_info->record.metric_id,
  };
  const uint64_t hash = iree_profile_model_metric_descriptor_hash(
      descriptor_info->record.source_id, descriptor_info->record.metric_id);
  iree_host_size_t existing_index = 0;
  if (iree_profile_index_find(&model->metric_descriptor_index, hash,
                              iree_profile_model_metric_descriptor_matches,
                              &lookup, &existing_index)) {
    return iree_ok_status();
  }

  if (model->metric_descriptor_count + 1 > model->metric_descriptor_capacity) {
    IREE_RETURN_IF_ERROR(iree_allocator_grow_array(
        model->host_allocator,
        iree_max((iree_host_size_t)16, model->metric_descriptor_count + 1),
        sizeof(model->metric_descriptors[0]),
        &model->metric_descriptor_capacity,
        (void**)&model->metric_descriptors));
  }
  const iree_host_size_t descriptor_index = model->metric_descriptor_count;
  IREE_RETURN_IF_ERROR(
      iree_profile_index_insert(&model->metric_descriptor_index,
                                model->host_allocator, hash, descriptor_index));
  model->metric_descriptors[descriptor_index] = *descriptor_info;
  ++model->metric_descriptor_count;
  return iree_ok_status();
}

static iree_status_t iree_profile_model_process_queue_records(
    iree_profile_model_t* model, const iree_hal_profile_file_record_t* record) {
  iree_profile_typed_record_iterator_t iterator;
  iree_profile_typed_record_iterator_initialize(
      record, sizeof(iree_hal_profile_queue_record_t), &iterator);
  iree_status_t status = iree_ok_status();
  while (iree_status_is_ok(status)) {
    iree_profile_typed_record_t typed_record;
    bool has_record = false;
    status = iree_profile_typed_record_iterator_next(&iterator, &typed_record,
                                                     &has_record);
    if (!iree_status_is_ok(status) || !has_record) break;

    iree_hal_profile_queue_record_t record_value;
    memcpy(&record_value, typed_record.contents.data, sizeof(record_value));
    status = iree_profile_model_append_queue(model, &record_value);
  }
  return status;
}

static iree_status_t iree_profile_model_process_metric_source_records(
    iree_profile_model_t* model, const iree_hal_profile_file_record_t* record) {
  iree_profile_typed_record_iterator_t iterator;
  iree_profile_typed_record_iterator_initialize(
      record, sizeof(iree_hal_profile_device_metric_source_record_t),
      &iterator);
  iree_status_t status = iree_ok_status();
  while (iree_status_is_ok(status)) {
    iree_profile_typed_record_t typed_record;
    bool has_record = false;
    status = iree_profile_typed_record_iterator_next(&iterator, &typed_record,
                                                     &has_record);
    if (!iree_status_is_ok(status) || !has_record) break;

    iree_hal_profile_device_metric_source_record_t record_value;
    memcpy(&record_value, typed_record.contents.data, sizeof(record_value));
    if ((iree_host_size_t)record_value.name_length !=
        typed_record.inline_payload.data_length) {
      status = iree_make_status(IREE_STATUS_DATA_LOSS,
                                "device metric source name length is "
                                "inconsistent with record length");
    }
    if (iree_status_is_ok(status)) {
      iree_profile_model_metric_source_t source_info = {
          .record = record_value,
          .name = iree_make_string_view(
              (const char*)typed_record.inline_payload.data,
              typed_record.inline_payload.data_length),
      };
      status = iree_profile_model_append_metric_source(model, &source_info);
    }
  }
  return status;
}

static iree_status_t iree_profile_model_process_metric_descriptor_records(
    iree_profile_model_t* model, const iree_hal_profile_file_record_t* record) {
  iree_profile_typed_record_iterator_t iterator;
  iree_profile_typed_record_iterator_initialize(
      record, sizeof(iree_hal_profile_device_metric_descriptor_record_t),
      &iterator);
  iree_status_t status = iree_ok_status();
  while (iree_status_is_ok(status)) {
    iree_profile_typed_record_t typed_record;
    bool has_record = false;
    status = iree_profile_typed_record_iterator_next(&iterator, &typed_record,
                                                     &has_record);
    if (!iree_status_is_ok(status) || !has_record) break;

    iree_hal_profile_device_metric_descriptor_record_t record_value;
    memcpy(&record_value, typed_record.contents.data, sizeof(record_value));
    iree_host_size_t trailing_length = 0;
    if (!iree_host_size_checked_add(record_value.name_length,
                                    record_value.description_length,
                                    &trailing_length) ||
        trailing_length != typed_record.inline_payload.data_length) {
      status = iree_make_status(
          IREE_STATUS_DATA_LOSS,
          "device metric descriptor string lengths are inconsistent with "
          "record length");
    }
    if (iree_status_is_ok(status)) {
      const char* string_base = (const char*)typed_record.inline_payload.data;
      iree_profile_model_metric_descriptor_t descriptor_info = {
          .record = record_value,
          .name = iree_make_string_view(string_base, record_value.name_length),
          .description =
              iree_make_string_view(string_base + record_value.name_length,
                                    record_value.description_length),
      };
      status =
          iree_profile_model_append_metric_descriptor(model, &descriptor_info);
    }
  }
  return status;
}

static iree_status_t iree_profile_model_process_executable_records(
    iree_profile_model_t* model, const iree_hal_profile_file_record_t* record) {
  iree_profile_typed_record_iterator_t iterator;
  iree_profile_typed_record_iterator_initialize(
      record, sizeof(iree_hal_profile_executable_record_t), &iterator);
  iree_status_t status = iree_ok_status();
  while (iree_status_is_ok(status)) {
    iree_profile_typed_record_t typed_record;
    bool has_record = false;
    status = iree_profile_typed_record_iterator_next(&iterator, &typed_record,
                                                     &has_record);
    if (!iree_status_is_ok(status) || !has_record) break;

    iree_hal_profile_executable_record_t record_value;
    memcpy(&record_value, typed_record.contents.data, sizeof(record_value));
    status = iree_profile_model_append_executable(model, &record_value);
  }
  return status;
}

static iree_status_t iree_profile_model_process_function_records(
    iree_profile_model_t* model, const iree_hal_profile_file_record_t* record) {
  iree_profile_typed_record_iterator_t iterator;
  iree_profile_typed_record_iterator_initialize(
      record, sizeof(iree_hal_profile_executable_function_record_t), &iterator);
  iree_status_t status = iree_ok_status();
  while (iree_status_is_ok(status)) {
    iree_profile_typed_record_t typed_record;
    bool has_record = false;
    status = iree_profile_typed_record_iterator_next(&iterator, &typed_record,
                                                     &has_record);
    if (!iree_status_is_ok(status) || !has_record) break;

    iree_hal_profile_executable_function_record_t record_value;
    memcpy(&record_value, typed_record.contents.data, sizeof(record_value));
    if ((iree_host_size_t)record_value.name_length !=
        typed_record.inline_payload.data_length) {
      status =
          iree_make_status(IREE_STATUS_DATA_LOSS,
                           "executable function name length is inconsistent");
    }
    if (iree_status_is_ok(status)) {
      iree_profile_model_function_t function_info = {
          .executable_id = record_value.executable_id,
          .flags = record_value.flags,
          .function_ordinal = record_value.function_ordinal,
          .constant_count = record_value.constant_count,
          .binding_count = record_value.binding_count,
          .parameter_count = record_value.parameter_count,
          .workgroup_size = {record_value.workgroup_size[0],
                             record_value.workgroup_size[1],
                             record_value.workgroup_size[2]},
          .function_hash = {record_value.function_hash[0],
                            record_value.function_hash[1]},
          .name = iree_make_string_view(
              (const char*)typed_record.inline_payload.data,
              typed_record.inline_payload.data_length),
      };
      status = iree_profile_model_append_function(model, &function_info);
    }
  }
  return status;
}

static iree_status_t iree_profile_model_process_command_buffer_records(
    iree_profile_model_t* model, const iree_hal_profile_file_record_t* record) {
  iree_profile_typed_record_iterator_t iterator;
  iree_profile_typed_record_iterator_initialize(
      record, sizeof(iree_hal_profile_command_buffer_record_t), &iterator);
  iree_status_t status = iree_ok_status();
  while (iree_status_is_ok(status)) {
    iree_profile_typed_record_t typed_record;
    bool has_record = false;
    status = iree_profile_typed_record_iterator_next(&iterator, &typed_record,
                                                     &has_record);
    if (!iree_status_is_ok(status) || !has_record) break;

    iree_hal_profile_command_buffer_record_t record_value;
    memcpy(&record_value, typed_record.contents.data, sizeof(record_value));
    status = iree_profile_model_append_command_buffer(model, &record_value);
  }
  return status;
}

static iree_status_t iree_profile_model_process_command_operation_records(
    iree_profile_model_t* model, const iree_hal_profile_file_record_t* record) {
  iree_profile_typed_record_iterator_t iterator;
  iree_profile_typed_record_iterator_initialize(
      record, sizeof(iree_hal_profile_command_operation_record_t), &iterator);
  iree_status_t status = iree_ok_status();
  while (iree_status_is_ok(status)) {
    iree_profile_typed_record_t typed_record;
    bool has_record = false;
    status = iree_profile_typed_record_iterator_next(&iterator, &typed_record,
                                                     &has_record);
    if (!iree_status_is_ok(status) || !has_record) break;

    iree_hal_profile_command_operation_record_t record_value;
    memcpy(&record_value, typed_record.contents.data, sizeof(record_value));
    status = iree_profile_model_validate_command_operation(&record_value);
    if (iree_status_is_ok(status)) {
      status =
          iree_profile_model_append_command_operation(model, &record_value);
    }
  }
  return status;
}

static iree_status_t iree_profile_model_process_clock_records(
    iree_profile_model_t* model, const iree_hal_profile_file_record_t* record) {
  iree_profile_typed_record_iterator_t iterator;
  iree_profile_typed_record_iterator_initialize(
      record, sizeof(iree_hal_profile_clock_correlation_record_t), &iterator);
  iree_status_t status = iree_ok_status();
  while (iree_status_is_ok(status)) {
    iree_profile_typed_record_t typed_record;
    bool has_record = false;
    status = iree_profile_typed_record_iterator_next(&iterator, &typed_record,
                                                     &has_record);
    if (!iree_status_is_ok(status) || !has_record) break;

    iree_hal_profile_clock_correlation_record_t clock_record;
    memcpy(&clock_record, typed_record.contents.data, sizeof(clock_record));

    iree_profile_model_device_t* device = NULL;
    status = iree_profile_model_ensure_device(
        model, clock_record.physical_device_ordinal, &device);
    if (iree_status_is_ok(status)) {
      iree_profile_model_record_clock_sample(device, &clock_record);
    }
  }
  return status;
}

typedef iree_status_t (*iree_profile_model_chunk_processor_fn_t)(
    iree_profile_model_t* model, const iree_hal_profile_file_record_t* record);

typedef struct iree_profile_model_chunk_route_t {
  // Profile chunk content type handled by this route.
  iree_string_view_t content_type;
  // Processor used to add the chunk's metadata rows to the model.
  iree_profile_model_chunk_processor_fn_t process;
} iree_profile_model_chunk_route_t;

iree_status_t iree_profile_model_process_metadata_record(
    iree_profile_model_t* model, const iree_hal_profile_file_record_t* record) {
  if (record->header.record_type != IREE_HAL_PROFILE_FILE_RECORD_TYPE_CHUNK) {
    return iree_ok_status();
  }
  const iree_profile_model_chunk_route_t routes[] = {
      {IREE_HAL_PROFILE_CONTENT_TYPE_QUEUES,
       iree_profile_model_process_queue_records},
      {IREE_HAL_PROFILE_CONTENT_TYPE_EXECUTABLES,
       iree_profile_model_process_executable_records},
      {IREE_HAL_PROFILE_CONTENT_TYPE_EXECUTABLE_FUNCTIONS,
       iree_profile_model_process_function_records},
      {IREE_HAL_PROFILE_CONTENT_TYPE_COMMAND_BUFFERS,
       iree_profile_model_process_command_buffer_records},
      {IREE_HAL_PROFILE_CONTENT_TYPE_COMMAND_OPERATIONS,
       iree_profile_model_process_command_operation_records},
      {IREE_HAL_PROFILE_CONTENT_TYPE_CLOCK_CORRELATIONS,
       iree_profile_model_process_clock_records},
      {IREE_HAL_PROFILE_CONTENT_TYPE_DEVICE_METRIC_SOURCES,
       iree_profile_model_process_metric_source_records},
      {IREE_HAL_PROFILE_CONTENT_TYPE_DEVICE_METRIC_DESCRIPTORS,
       iree_profile_model_process_metric_descriptor_records},
  };
  for (iree_host_size_t i = 0; i < IREE_ARRAYSIZE(routes); ++i) {
    if (iree_string_view_equal(record->content_type, routes[i].content_type)) {
      return routes[i].process(model, record);
    }
  }
  return iree_ok_status();
}

const char* iree_profile_command_operation_type_name(
    iree_hal_profile_command_operation_type_t type) {
  switch (type) {
    case IREE_HAL_PROFILE_COMMAND_OPERATION_TYPE_BARRIER:
      return "barrier";
    case IREE_HAL_PROFILE_COMMAND_OPERATION_TYPE_DISPATCH:
      return "dispatch";
    case IREE_HAL_PROFILE_COMMAND_OPERATION_TYPE_FILL:
      return "fill";
    case IREE_HAL_PROFILE_COMMAND_OPERATION_TYPE_COPY:
      return "copy";
    case IREE_HAL_PROFILE_COMMAND_OPERATION_TYPE_UPDATE:
      return "update";
    case IREE_HAL_PROFILE_COMMAND_OPERATION_TYPE_PROFILE_MARKER:
      return "profile_marker";
    case IREE_HAL_PROFILE_COMMAND_OPERATION_TYPE_BRANCH:
      return "branch";
    case IREE_HAL_PROFILE_COMMAND_OPERATION_TYPE_COND_BRANCH:
      return "cond_branch";
    case IREE_HAL_PROFILE_COMMAND_OPERATION_TYPE_RETURN:
      return "return";
    default:
      return "unknown";
  }
}

const char* iree_profile_queue_event_type_name(
    iree_hal_profile_queue_event_type_t type) {
  switch (type) {
    case IREE_HAL_PROFILE_QUEUE_EVENT_TYPE_BARRIER:
      return "barrier";
    case IREE_HAL_PROFILE_QUEUE_EVENT_TYPE_DISPATCH:
      return "dispatch";
    case IREE_HAL_PROFILE_QUEUE_EVENT_TYPE_EXECUTE:
      return "execute";
    case IREE_HAL_PROFILE_QUEUE_EVENT_TYPE_COPY:
      return "copy";
    case IREE_HAL_PROFILE_QUEUE_EVENT_TYPE_FILL:
      return "fill";
    case IREE_HAL_PROFILE_QUEUE_EVENT_TYPE_UPDATE:
      return "update";
    case IREE_HAL_PROFILE_QUEUE_EVENT_TYPE_READ:
      return "read";
    case IREE_HAL_PROFILE_QUEUE_EVENT_TYPE_WRITE:
      return "write";
    case IREE_HAL_PROFILE_QUEUE_EVENT_TYPE_ALLOCA:
      return "alloca";
    case IREE_HAL_PROFILE_QUEUE_EVENT_TYPE_DEALLOCA:
      return "dealloca";
    case IREE_HAL_PROFILE_QUEUE_EVENT_TYPE_HOST_CALL:
      return "host_call";
    default:
      return "unknown";
  }
}

const char* iree_profile_queue_dependency_strategy_name(
    iree_hal_profile_queue_dependency_strategy_t strategy) {
  switch (strategy) {
    case IREE_HAL_PROFILE_QUEUE_DEPENDENCY_STRATEGY_NONE:
      return "none";
    case IREE_HAL_PROFILE_QUEUE_DEPENDENCY_STRATEGY_INLINE:
      return "inline";
    case IREE_HAL_PROFILE_QUEUE_DEPENDENCY_STRATEGY_DEVICE_BARRIER:
      return "device_barrier";
    case IREE_HAL_PROFILE_QUEUE_DEPENDENCY_STRATEGY_SOFTWARE_DEFER:
      return "software_defer";
    default:
      return "unknown";
  }
}

const char* iree_profile_event_relationship_type_name(
    iree_hal_profile_event_relationship_type_t type) {
  switch (type) {
    case IREE_HAL_PROFILE_EVENT_RELATIONSHIP_TYPE_QUEUE_SUBMISSION_DISPATCH:
      return "queue_submission_dispatch";
    case IREE_HAL_PROFILE_EVENT_RELATIONSHIP_TYPE_QUEUE_SUBMISSION_QUEUE_DEVICE_EVENT:
      return "queue_submission_queue_device_event";
    case IREE_HAL_PROFILE_EVENT_RELATIONSHIP_TYPE_QUEUE_EVENT_HOST_EXECUTION_EVENT:
      return "queue_event_host_execution_event";
    case IREE_HAL_PROFILE_EVENT_RELATIONSHIP_TYPE_QUEUE_SUBMISSION_HOST_EXECUTION_EVENT:
      return "queue_submission_host_execution_event";
    default:
      return "unknown";
  }
}

const char* iree_profile_event_endpoint_type_name(
    iree_hal_profile_event_endpoint_type_t type) {
  switch (type) {
    case IREE_HAL_PROFILE_EVENT_ENDPOINT_TYPE_QUEUE_SUBMISSION:
      return "queue_submission";
    case IREE_HAL_PROFILE_EVENT_ENDPOINT_TYPE_QUEUE_EVENT:
      return "queue_event";
    case IREE_HAL_PROFILE_EVENT_ENDPOINT_TYPE_DISPATCH_EVENT:
      return "dispatch_event";
    case IREE_HAL_PROFILE_EVENT_ENDPOINT_TYPE_COMMAND_OPERATION:
      return "command_operation";
    case IREE_HAL_PROFILE_EVENT_ENDPOINT_TYPE_MEMORY_EVENT:
      return "memory_event";
    case IREE_HAL_PROFILE_EVENT_ENDPOINT_TYPE_ARTIFACT:
      return "artifact";
    case IREE_HAL_PROFILE_EVENT_ENDPOINT_TYPE_QUEUE_DEVICE_EVENT:
      return "queue_device_event";
    case IREE_HAL_PROFILE_EVENT_ENDPOINT_TYPE_HOST_EXECUTION_EVENT:
      return "host_execution_event";
    default:
      return "unknown";
  }
}

double iree_profile_model_span_ticks(uint64_t earliest_start_tick,
                                     uint64_t latest_end_tick) {
  if (earliest_start_tick == UINT64_MAX || latest_end_tick == 0 ||
      latest_end_tick < earliest_start_tick) {
    return 0.0;
  }
  return (double)(latest_end_tick - earliest_start_tick);
}

iree_status_t iree_profile_model_resolve_command_operation_key(
    const iree_profile_model_t* model,
    const iree_hal_profile_command_operation_record_t* operation,
    char* numeric_buffer, iree_host_size_t numeric_buffer_capacity,
    iree_string_view_t* out_key) {
  *out_key = iree_make_cstring_view(
      iree_profile_command_operation_type_name(operation->type));
  if (operation->type != IREE_HAL_PROFILE_COMMAND_OPERATION_TYPE_DISPATCH ||
      operation->executable_id == 0 ||
      operation->function_ordinal == UINT32_MAX) {
    return iree_ok_status();
  }

  const iree_profile_model_command_buffer_t* command_buffer =
      iree_profile_model_find_command_buffer(model,
                                             operation->command_buffer_id);
  if (!command_buffer) {
    return iree_make_status(
        IREE_STATUS_DATA_LOSS,
        "command operation references missing command-buffer metadata "
        "command_buffer=%" PRIu64 " command_index=%u",
        operation->command_buffer_id, operation->command_index);
  }
  return iree_profile_model_resolve_dispatch_key(
      model, command_buffer->record.physical_device_ordinal,
      operation->executable_id, operation->function_ordinal, numeric_buffer,
      numeric_buffer_capacity, out_key);
}
