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

static const char kIreeProfileUsage[] =
    "Inspects IREE HAL profile bundles.\n"
    "\n"
    "A .ireeprof bundle is a HAL-native profiling artifact. It contains\n"
    "session metadata, physical-device and queue metadata, "
    "executable/function\n"
    "metadata, command-buffer metadata and operations, clock-correlation\n"
    "samples, host queue-operation events, device-timestamped dispatch "
    "events,\n"
    "optional memory lifecycle events, and optional hardware counter samples. "
    "The projection commands below let you enter "
    "from\n"
    "the object you care about and then cross-reference ids in JSONL.\n"
    "\n"
    "Usage:\n"
    "  iree-profile summary [--format=text|jsonl] <file.ireeprof>\n"
    "  iree-profile statistics [--format=text|jsonl] <file.ireeprof>\n"
    "  iree-profile att [--format=text|jsonl] [--filter=pattern]\n"
    "      [--id=trace_or_dispatch_event_id]\n"
    "      [--rocm_library_path=/opt/rocm/lib] <file.ireeprof>\n"
    "      (when built with AMDGPU profiling support)\n"
    "  iree-profile explain [--format=text|jsonl] [--filter=pattern]\n"
    "      [--id=dispatch_event_id] <file.ireeprof>\n"
    "  iree-profile executable [--format=text|jsonl] [--id=executable_id]\n"
    "      [--filter=pattern] [--dispatch_events] <file.ireeprof>\n"
    "  iree-profile dispatch [--format=text|jsonl] [--id=event_id]\n"
    "      [--filter=pattern] [--dispatch_events] <file.ireeprof>\n"
    "  iree-profile command [--format=text|jsonl] [--id=command_buffer_id]\n"
    "      [--filter=pattern] [--dispatch_events] <file.ireeprof>\n"
    "  iree-profile counter [--format=text|jsonl] "
    "[--id=sample_or_related_id]\n"
    "      [--filter=pattern] [--counter_samples] <file.ireeprof>\n"
    "  iree-profile memory [--format=text|jsonl] "
    "[--id=event_or_allocation_id]\n"
    "      [--filter=event_type_pattern] <file.ireeprof>\n"
    "  iree-profile queue [--format=text|jsonl] [--id=submission_id]\n"
    "      [--filter=pattern] [--dispatch_events] <file.ireeprof>\n"
    "  iree-profile export --format=ireeperf-jsonl [--output=path|-]\n"
    "      <file.ireeprof>\n"
    "  iree-profile cat [--format=text|jsonl] <file.ireeprof>\n"
    "  iree-profile --agents_md\n"
    "\n"
    "Commands:\n"
    "  summary      Bundle health, metadata counts, clock fit, and per-device\n"
    "               dispatch timing totals.\n"
    "  statistics   Aggregate statistics reduced directly from raw profile\n"
    "               chunks. This mirrors the lightweight runtime statistics\n"
    "               sink and is useful for scripts that want one compact\n"
    "               row stream.\n"
    "  att          Optional AMDGPU ATT/SQTT trace decode when built with\n"
    "               AMDGPU profiling support. Uses ROCm's\n"
    "               trace decoder and annotates decoded PCs with embedded\n"
    "               code-object disassembly.\n"
    "  explain      Opinionated bottleneck summary: device spans, queue busy\n"
    "               intervals, top functions/dispatches, transfer/memory "
    "totals,"
    "\n"
    "               and evidence-backed hints.\n"
    "  executable   Executable/function catalog joined with dispatch timing. "
    "Use\n"
    "               this when optimizing a specific kernel name or function. "
    "Includes\n"
    "               code_object_hash and function_hash when producers provide "
    "them.\n"
    "  dispatch     Per-function dispatch timing aggregates, or individual "
    "events\n"
    "               with --dispatch_events --format=jsonl.\n"
    "  command      Recorded command-buffer operations, metadata, and "
    "per-execution\n"
    "               dispatch spans. Use this to inspect copy/fill/update/"
    "barrier\n"
    "               structure and drill into dispatch timings.\n"
    "  queue        Queue operation events and dispatch-derived submission "
    "spans.\n"
    "  counter      Hardware counter metadata and sample aggregates joined "
    "to\n"
    "               dispatch/function/command/queue ids. Use "
    "--counter_samples\n"
    "               with JSONL for per-dispatch sample rows.\n"
    "  memory       Memory lifecycle events and high-water summaries for "
    "slabs,\n"
    "               pools, async queue allocations, and HAL buffers. "
    "Live-at-end counts\n"
    "               are capture-window state, not required leaks.\n"
    "  export       Decoded tooling interchange export. The first supported\n"
    "               format is schema-versioned ireeperf-jsonl.\n"
    "  cat          Raw bundle record dump for format archaeology/debugging.\n"
    "\n"
    "Important flags:\n"
    "  --format=FORMAT         text/jsonl for reports; ireeperf-jsonl for\n"
    "                          function interchange.\n"
    "  --filter=pattern        Wildcard over dispatch/function keys, command "
    "op\n"
    "                          names/keys, or queue/memory event type names,\n"
    "                          such as '*softmax*' or 'copy'.\n"
    "  --id=N                  att: trace_id or dispatch event_id;\n"
    "                          dispatch: event_id; executable: executable_id;\n"
    "                          command: command_buffer_id; memory: "
    "event_id or\n"
    "                          allocation_id; queue: submission_id; "
    "counter:\n"
    "                          sample_id, dispatch_event_id, "
    "submission_id, or\n"
    "                          command_buffer_id; explain:\n"
    "                          dispatch event id.\n"
    "  --dispatch_events       Stream individual dispatch_event rows instead "
    "of\n"
    "                          only aggregate projection rows. Requires "
    "JSONL.\n"
    "  --counter_samples       Stream individual counter_sample rows instead "
    "of\n"
    "                          only aggregate counter rows. Requires JSONL.\n"
    "  --output=path|-         Export destination path, or `-` for stdout.\n"
    "  --rocm_library_path=path ROCm library directory or exact dynamic "
    "library\n"
    "                          path used by `att`. Overrides\n"
    "                          IREE_HAL_AMDGPU_LIBAQLPROFILE_PATH and\n"
    "                          IREE_HAL_AMDGPU_LIBHSA_PATH; if all are "
    "unset,\n"
    "                          uses the system dynamic library search path.\n"
    "  --agents_md             Print a Markdown guide optimized for "
    "AGENTS.md.\n"
    "\n"
    "JSONL contracts:\n"
    "  Projection commands with --format=jsonl emit command-local report rows\n"
    "  keyed by `type`. They are meant for jq, agents, and human drilldown.\n"
    "  function --format=ireeperf-jsonl emits schema-versioned interchange "
    "rows\n"
    "  keyed by `record_type`. Use that stream for downstream adapters,\n"
    "  telemetry import, and long-lived tooling.\n"
    "\n"
    "Capture examples:\n"
    "  iree-benchmark-module --device=amdgpu --module=model.vmfb \\\n"
    "      --function=main --benchmark_min_time=20x \\\n"
    "      --device_profiling_mode=queue-events,device-queue-events \\\n"
    "      --device_profiling_output=/tmp/model.ireeprof\n"
    "\n"
    "Embedding examples:\n"
    "  Use iree_hal_profile_file_sink_create to create a file-backed sink,\n"
    "  set iree_hal_device_profiling_options_t::sink, then call\n"
    "  iree_hal_device_profiling_begin/end around the work to capture.\n"
    "\n"
    "Analysis examples:\n"
    "  iree-profile summary /tmp/model.ireeprof\n"
    "  iree-profile statistics --format=jsonl /tmp/model.ireeprof | \\\n"
    "      jq 'select(.type==\"statistics_row\") | {row_type,event_name,"
    "total_duration_ns}'\n"
    "  iree-profile att --rocm_library_path=/opt/rocm/lib \\\n"
    "      --filter='*matmul*' /tmp/model.ireeprof\n"
    "  iree-profile explain /tmp/model.ireeprof\n"
    "  iree-profile explain --format=jsonl /tmp/model.ireeprof | \\\n"
    "      jq 'select(.type | startswith(\"explain_top_\"))'\n"
    "  iree-profile export --format=ireeperf-jsonl \\\n"
    "      --output=/tmp/model.ireeperf.jsonl /tmp/model.ireeprof\n"
    "  iree-profile executable --filter='*matmul*' /tmp/model.ireeprof\n"
    "  iree-profile dispatch --format=jsonl /tmp/model.ireeprof | \\\n"
    "      jq 'select(.type==\"dispatch_group\") | {key,avg_ns,count}'\n"
    "  iree-profile queue --format=jsonl /tmp/model.ireeprof | \\\n"
    "      jq 'select(.type==\"queue_event\" or "
    ".type==\"queue_device_event\" or .type==\"queue_submission\")'\n"
    "  iree-profile counter --format=jsonl /tmp/model.ireeprof | \\\n"
    "      jq 'select(.type==\"counter_group\") | "
    "{key,counter,avg,sum,samples}'\n"
    "  iree-profile counter --format=jsonl --counter_samples \\\n"
    "      /tmp/model.ireeprof | jq 'select(.type==\"counter_sample\")'\n"
    "  iree-profile memory --format=jsonl /tmp/model.ireeprof | \\\n"
    "      jq 'select(.type==\"memory_pool\") | \\\n"
    "          {kind,physical_device_ordinal,pool_id,"
    "lifecycle_high_water_bytes,waits}'\n"
    "  iree-profile command --id=1 --format=jsonl --dispatch_events \\\n"
    "      /tmp/model.ireeprof | jq 'select(.type==\"dispatch_event\")'\n"
    "\n"
    "Use `iree-profile --agents_md` for a Markdown playbook with JSONL record\n"
    "types and cross-reference recipes.\n";

void iree_profile_fprint_usage(FILE* file) {
  fputs("iree-profile\n", file);
  fputs(kIreeProfileUsage, file);
}

static void iree_profile_print_agent_markdown_capture(FILE* file) {
  fputs(
      "# IREE HAL Profiling With iree-profile\n"
      "\n"
      "Use `iree-profile` to inspect `.ireeprof` bundles emitted by IREE HAL\n"
      "device profiling. Prefer `--format=jsonl` for automated analysis: each\n"
      "line is one independent JSON record that can be filtered with `jq`.\n"
      "\n"
      "## Capture\n"
      "\n"
      "Capture from `iree-benchmark-module` with HAL queue profiling "
      "enabled:\n"
      "\n"
      "```bash\n"
      "iree-benchmark-module --device=amdgpu --module=model.vmfb \\\n"
      "  --function=main --benchmark_min_time=20x \\\n"
      "  --device_profiling_mode=queue-events,device-queue-events \\\n"
      "  --device_profiling_output=/tmp/model.ireeprof\n"
      "```\n"
      "\n"
      "Embedding applications can capture the same bundle through the HAL "
      "API:\n"
      "\n"
      "```c\n"
      "iree_io_file_handle_t* file_handle = NULL;\n"
      "iree_hal_profile_sink_t* sink = NULL;\n"
      "IREE_RETURN_IF_ERROR(iree_io_file_handle_create(\n"
      "    IREE_IO_FILE_MODE_WRITE | IREE_IO_FILE_MODE_SEQUENTIAL_SCAN,\n"
      "    IREE_SV(\"/tmp/model.ireeprof\"), 0, host_allocator, "
      "&file_handle));\n"
      "IREE_RETURN_IF_ERROR(iree_hal_profile_file_sink_create(\n"
      "    file_handle, host_allocator, &sink));\n"
      "iree_hal_device_profiling_options_t options = {0};\n"
      "options.data_families = IREE_HAL_DEVICE_PROFILING_DATA_QUEUE_EVENTS |\n"
      "                        "
      "IREE_HAL_DEVICE_PROFILING_DATA_DEVICE_QUEUE_EVENTS;\n"
      "options.sink = sink;\n"
      "IREE_RETURN_IF_ERROR(iree_hal_device_profiling_begin(device, "
      "&options));\n"
      "/* Run workload. */\n"
      "IREE_RETURN_IF_ERROR(iree_hal_device_profiling_end(device));\n"
      "iree_hal_profile_sink_release(sink);\n"
      "iree_io_file_handle_release(file_handle);\n"
      "```\n",
      file);
}

static void iree_profile_print_agent_markdown_commands(FILE* file) {
  fputs(
      "\n"
      "## Commands\n"
      "\n"
      "- `summary` checks bundle health, metadata counts, dispatch counts, "
      "clock\n"
      "  correlation, and per-device timing totals.\n"
      "- `statistics` reduces raw profile chunks through the same aggregate "
      "sink\n"
      "  used by lightweight runtime statistics. It emits compact row "
      "families\n"
      "  such as dispatch-by-function, queue operation, host execution, and "
      "memory\n"
      "  lifecycle aggregates without command-specific drilldown joins.\n"
      "- `att`, when built, decodes AMDGPU ATT/SQTT executable trace blobs "
      "with ROCm's\n"
      "  trace decoder and annotates decoded PCs using code objects embedded "
      "in\n"
      "  the `.ireeprof` bundle. Library lookup uses `--rocm_library_path` "
      "first,\n"
      "  then `IREE_HAL_AMDGPU_LIBAQLPROFILE_PATH`, then "
      "`IREE_HAL_AMDGPU_LIBHSA_PATH`,\n"
      "  then the normal system dynamic library search path.\n"
      "- `explain` gives an opinionated first-pass bottleneck view: visible "
      "device\n"
      "  spans, summed active dispatch time, merged per-queue busy intervals, "
      "top\n"
      "  functions, top individual dispatches, transfer totals, memory "
      "pressure,\n"
      "  and evidence-backed hints. Queue gap hints require queue event "
      "records;\n"
      "  do not infer queue dependency causes from dispatch record order.\n"
      "- `executable` lists executable/function metadata and joins function "
      "ids "
      "to\n"
      "  dispatch timing aggregates. Hash-capable producers include "
      "`code_object_hash`\n"
      "  on `executable` rows and `function_hash` on `executable_function` "
      "rows for\n"
      "  cross-run and external-tool correlation.\n"
      "- `dispatch` groups timings by executable function, or emits every "
      "dispatch\n"
      "  event with `--dispatch_events`.\n"
      "- `command` lists command-buffer metadata, static operation records, "
      "and\n"
      "  execution spans grouped by `command_buffer_id` and `submission_id`.\n"
      "- `queue` groups dispatch-derived queue submissions by physical queue "
      "and\n"
      "  `submission_id`, and emits host queue-operation events when the "
      "producer\n"
      "  provides them. Queue events include operation type, wait/signal "
      "counts,\n"
      "  dependency strategy, payload bytes, and related object ids. "
      "Queue-device\n"
      "  events add raw device start/end ticks for timed submissions such as "
      "alloca,\n"
      "  execute, and dealloca.\n"
      "- `counter` joins hardware counter metadata and samples to "
      "dispatch/function,\n"
      "  command-buffer, queue, and submission ids. Aggregate mode emits "
      "per-counter\n"
      "  statistics; `--counter_samples --format=jsonl` emits every matched "
      "sample\n"
      "  row with raw values for drilldown.\n"
      "- `memory` summarizes slab/provider events, pool reservation events, "
      "async\n"
      "  queue alloca/dealloca, and synchronous HAL buffer allocation/free "
      "high-water behavior. Externally-owned imports are reported separately "
      "from IREE-owned allocations. With `--format=jsonl` it emits individual\n"
      "  `memory_event` rows before the summary and then emits `memory_pool` "
      "and\n"
      "  `memory_allocation` rows for pool/provider and lifecycle drilldown. "
      "It\n"
      "  separates reserved pool bytes, materialized pool bytes, queue "
      "in-flight\n"
      "  bytes, synchronous buffer bytes, and pool free/budget snapshots when "
      "the\n"
      "  producer supplies them. Open lifecycles at profile end or\n"
      "  partial closes can be normal capture-window/filter artifacts.\n"
      "- `export --format=ireeperf-jsonl` decodes `.ireeprof` into a "
      "schema-versioned\n"
      "  tooling interchange stream. This is the stable boundary for "
      "downstream\n"
      "  adapters, telemetry upload, database import, and optional renderers "
      "such as Perfetto or "
      "SQLite.\n"
      "- `cat` dumps raw bundle records and is mainly for format debugging.\n",
      file);
}

static void iree_profile_print_agent_markdown_jsonl_contracts(FILE* file) {
  fputs(
      "\n"
      "## JSONL Contracts\n"
      "\n"
      "Projection commands with `--format=jsonl` emit command-local report "
      "rows.\n"
      "Those rows are keyed by `type` and are optimized for interactive "
      "`jq`,\n"
      "agent analysis, and command-specific drilldown. The field names are "
      "kept\n"
      "boring and consistent across commands where the same concept appears, "
      "but\n"
      "these rows are still CLI reports, not the decoded bundle "
      "interchange\n"
      "schema.\n"
      "\n"
      "`export --format=ireeperf-jsonl` is the schema-versioned interchange "
      "path.\n"
      "Every row has `schema_version` and `record_type`, and the first row "
      "is a\n"
      "`schema` record describing time domains and clock-fit behavior. "
      "Use this\n"
      "stream for downstream adapters, telemetry upload, database imports, "
      "and\n"
      "renderers such as Perfetto/SQLite/HTML. Branch on `record_type` for "
      "export\n"
      "streams and on `type` for command report streams; do not merge the "
      "two\n"
      "families as if they had one row namespace.\n",
      file);
}

static void iree_profile_print_agent_markdown_jsonl_record_types(FILE* file) {
  fputs(
      "\n"
      "## JSONL Record Types\n"
      "\n"
      "- `summary` emits `summary` and `device_summary`.\n"
      "- `statistics --format=jsonl` emits `statistics_summary` and "
      "`statistics_row`.\n"
      "- `att --format=jsonl` emits `att_trace` and `att_instruction`.\n"
      "- `explain --format=jsonl` emits `explain_summary`, "
      "`explain_device`,\n"
      "  `explain_queue`, `explain_top_function`, "
      "`explain_top_dispatch`,\n"
      "  `explain_queue_operations`, `explain_memory_pressure`, and "
      "`explain_hint`.\n"
      "- `executable --format=jsonl` emits `executable_summary`, "
      "`executable`,\n"
      "  `executable_function`, and `executable_function_dispatch_group`.\n"
      "- `dispatch --format=jsonl` emits `dispatch_summary` and "
      "`dispatch_group`.\n"
      "- `dispatch --format=jsonl --dispatch_events` emits `dispatch_event` "
      "rows\n"
      "  followed by `dispatch_summary`.\n"
      "- `command --format=jsonl` emits `command_summary`, `command_buffer`,\n"
      "  `command_operation`, and `command_execution`.\n"
      "- `queue --format=jsonl` emits `queue_summary`, `queue`,\n"
      "  `queue_submission`, `queue_device_event`, and `queue_event`.\n"
      "- `counter --format=jsonl` emits `counter_summary`, `counter_set`,\n"
      "  `counter`, and `counter_group`.\n"
      "- `counter --format=jsonl --counter_samples` additionally emits\n"
      "  `counter_sample` rows before the summary.\n"
      "- `memory --format=jsonl` emits `memory_event`, `memory_summary`,\n"
      "  `memory_device`, `memory_pool`, and `memory_allocation`.\n"
      "- `export --format=ireeperf-jsonl` emits records with "
      "`schema_version`\n"
      "  and `record_type`: `schema`, `session`, `device`, `queue`,\n"
      "  `executable`, `executable_function`, `command_buffer`,\n"
      "  `command_operation`, `clock_correlation`, `dispatch_event`,\n"
      "  `queue_event`, `queue_device_event`, `memory_event`, "
      "`event_relationship`, "
      "`counter_set`,\n"
      "  `counter`, `counter_sample`, and `diagnostic`.\n",
      file);
}

static void iree_profile_print_agent_markdown_cross_references(FILE* file) {
  fputs(
      "\n"
      "## Cross-Reference Keys\n"
      "\n"
      "- `dispatch_event.event_id` uniquely identifies one dispatch event in "
      "the\n"
      "  profile stream. Use `dispatch --id=<event_id>` to isolate it.\n"
      "- `dispatch_event.submission_id` joins dispatches to "
      "`queue_submission`.\n"
      "  Use `queue --id=<submission_id>` or "
      "`queue --id=<submission_id> --dispatch_events` to inspect a "
      "submission.\n"
      "- `dispatch_event.command_buffer_id` joins dispatches to "
      "`command_buffer`\n"
      "  and `command_execution`. Use `command --id=<command_buffer_id>`.\n"
      "- `dispatch_event.command_index` is the ordinal of the dispatch within "
      "the\n"
      "  command buffer when the event came from reusable command-buffer "
      "replay.\n"
      "- `command_operation.command_buffer_id` and "
      "`command_operation.command_index`\n"
      "  join recorded command-buffer operations to dispatch events. "
      "`op` and\n"
      "  `key` can be filtered with `command --filter=<pattern>`.\n"
      "  `block_structure` means producer-local block ordinals are present;\n"
      "  linear or opaque producers leave those fields absent.\n"
      "- `dispatch_event.executable_id` plus `dispatch_event.function_ordinal` "
      "joins\n"
      "  to `executable_function`. The `key` field is the function name when "
      "present,\n"
      "  otherwise a stable numeric fallback.\n"
      "- `physical_device_ordinal`, `queue_ordinal`, and `stream_id` identify "
      "the\n"
      "  producing physical queue within the session.\n"
      "- `queue_event.submission_id` joins queue operation metadata to "
      "`queue_submission`\n"
      "  and dispatch events. `queue_event.command_buffer_id` joins execute "
      "events\n"
      "  to `command_buffer`; `queue_event.allocation_id` joins alloca/"
      "dealloca\n"
      "  events to memory events when memory lifecycle capture is enabled.\n"
      "- `queue_event.host_time_ns` is submission time; "
      "`queue_event.ready_host_time_ns`\n"
      "  is when the operation became eligible to execute, when known.\n"
      "- `queue_device_event.submission_id` joins host queue operations to "
      "device\n"
      "  start/end ticks in the same submission. For reusable command-buffer "
      "execute,\n"
      "  `queue_device_event.command_buffer_id` identifies the command buffer "
      "whose\n"
      "  whole-submission device span encloses its dispatch events. "
      "`queue_device_event`\n"
      "  records can be present even when `queue_event` host records were not "
      "requested.\n"
      "- `counter_sample.sample_id` identifies one producer sample. "
      "`counter_sample.scope`\n"
      "  classifies whether the values describe a dispatch, command "
      "operation, or\n"
      "  physical-device time range. When "
      "`counter_sample.device_tick_range_present`\n"
      "  is true, `start_tick` and `end_tick` are raw device timestamps in "
      "the\n"
      "  `device_tick` time domain.\n"
      "- "
      "`counter_sample.dispatch_event_id`, `submission_id`, "
      "`command_buffer_id`,\n"
      "  `command_index`, `executable_id`, `function_ordinal`, "
      "`physical_device_ordinal`,\n"
      "  `queue_ordinal`, and `stream_id` join counter data to the matching\n"
      "  dispatch, command, executable, and queue projections. "
      "`counter_set_id`\n"
      "  plus `counter_ordinal` joins samples to counter metadata.\n"
      "- `memory_event.submission_id` joins queue alloca/dealloca events to "
      "queue\n"
      "  submissions when nonzero. `memory_event.allocation_id` is the "
      "producer-defined session-local lifecycle id joining related memory "
      "events.\n"
      "  `memory_event.backing_id` and `memory_event.pool_id` are producer "
      "implementation identifiers and may be raw addresses or provider "
      "handles. "
      "`memory_pool`\n"
      "  groups by producer pool/provider id and memory type. "
      "`memory_allocation`\n"
      "  groups capture-window lifecycle rows by kind, allocation id, and "
      "pool id.\n"
      "  `memory_device`\n"
      "  live/high-water fields are computed over the capture window, so "
      "they\n"
      "  need not balance to zero if buffers outlive profiling.\n",
      file);
}

static void iree_profile_print_agent_markdown_recipes(FILE* file) {
  fputs(
      "\n"
      "## Recipes\n"
      "\n"
      "Find the slowest kernel/function groups:\n"
      "\n"
      "```bash\n"
      "iree-profile explain --format=jsonl /tmp/model.ireeprof | \\\n"
      "  jq 'select(.type==\"explain_top_function\") | \\\n"
      "      {rank,key,total_ns,avg_ns,max_ns,dispatches}'\n"
      "```\n"
      "\n"
      "Find the slowest individual dispatches:\n"
      "\n"
      "```bash\n"
      "iree-profile explain --format=jsonl /tmp/model.ireeprof | \\\n"
      "  jq 'select(.type==\"explain_top_dispatch\") | \\\n"
      "      {rank,event_id,submission_id,key,duration_ns}'\n"
      "```\n"
      "\n"
      "Build a full custom function table:\n"
      "\n"
      "```bash\n"
      "iree-profile dispatch --format=jsonl /tmp/model.ireeprof | \\\n"
      "  jq -s 'map(select(.type==\"dispatch_group\")) | \\\n"
      "         sort_by(.total_ns) | reverse[:20] | \\\n"
      "         map({key,count,avg_ns,total_ns,max_ns})'\n"
      "```\n"
      "\n"
      "Export the full decoded interchange stream for downstream tools:\n"
      "\n"
      "```bash\n"
      "iree-profile export --format=ireeperf-jsonl \\\n"
      "  --output=/tmp/model.ireeperf.jsonl /tmp/model.ireeprof\n"
      "jq 'select(.record_type==\"dispatch_event\") | \\\n"
      "    {event_id,key,duration_ns,duration_time_domain,"
      "derived_time_available,derived_time_domain}' \\\n"
      "  /tmp/model.ireeperf.jsonl\n"
      "```\n"
      "\n"
      "Show every dispatch of one kernel name pattern:\n"
      "\n"
      "```bash\n"
      "iree-profile dispatch --format=jsonl --dispatch_events \\\n"
      "  --filter='*softmax*' /tmp/model.ireeprof | \\\n"
      "  jq 'select(.type==\"dispatch_event\") | \\\n"
      "      {event_id,submission_id,command_buffer_id,command_index,key,"
      "duration_ns}'\n"
      "```\n"
      "\n"
      "Decode ATT traces for one kernel name pattern:\n"
      "\n"
      "```bash\n"
      "iree-profile att --format=jsonl --filter='*matmul*' \\\n"
      "  /tmp/model.ireeprof | \\\n"
      "  jq 'select(.type==\"att_instruction\") | \\\n"
      "      {trace_id,pc,category,hits,stall,instruction}'\n"
      "```\n"
      "\n"
      "Find expensive command-buffer executions:\n"
      "\n"
      "```bash\n"
      "iree-profile command --format=jsonl /tmp/model.ireeprof | \\\n"
      "  jq -s 'map(select(.type==\"command_execution\")) | \\\n"
      "         sort_by(.span_ns) | reverse[:20] | \\\n"
      "         map({command_buffer_id,submission_id,span_ns,dispatches})'\n"
      "```\n"
      "\n"
      "Drill into all dispatches for command buffer 1:\n"
      "\n"
      "```bash\n"
      "iree-profile command --id=1 --format=jsonl --dispatch_events \\\n"
      "  /tmp/model.ireeprof | jq 'select(.type==\"dispatch_event\")'\n"
      "```\n"
      "\n"
      "Show the recorded operation stream for command buffer 1:\n"
      "\n"
      "```bash\n"
      "iree-profile command --id=1 --format=jsonl /tmp/model.ireeprof | \\\n"
      "  jq 'select(.type==\"command_operation\") | \\\n"
      "      {command_index,op,key,length,workgroup_count}'\n"
      "```\n"
      "\n"
      "Find queue submissions with long device spans:\n"
      "\n"
      "```bash\n"
      "iree-profile queue --format=jsonl /tmp/model.ireeprof | \\\n"
      "  jq -s 'map(select(.type==\"queue_submission\")) | \\\n"
      "         sort_by(.span_ns) | reverse[:20] | \\\n"
      "         map({submission_id,physical_device_ordinal,queue_ordinal,"
      "span_ns,dispatches})'\n"
      "```\n"
      "\n"
      "Drill from queue submission 4 to its dispatch events:\n"
      "\n"
      "```bash\n"
      "iree-profile queue --id=4 --format=jsonl --dispatch_events \\\n"
      "  /tmp/model.ireeprof | jq 'select(.type==\"dispatch_event\")'\n"
      "```\n"
      "\n"
      "Summarize hardware counters by kernel/function:\n"
      "\n"
      "```bash\n"
      "iree-profile counter --format=jsonl /tmp/model.ireeprof | \\\n"
      "  jq -s 'map(select(.type==\"counter_group\")) | \\\n"
      "         sort_by(.sum) | reverse[:20] | \\\n"
      "         map({key,counter,samples,avg,sum})'\n"
      "```\n"
      "\n"
      "Capture low-disturbance PMC ranges without dispatch attribution:\n"
      "\n"
      "```bash\n"
      "iree-benchmark-module --device=amdgpu --module=model.vmfb \\\n"
      "  --function=main --benchmark_min_time=20x \\\n"
      "  --device_profiling_mode=counter-ranges \\\n"
      "  --device_profiling_counter=SQ_WAVES \\\n"
      "  --device_profiling_output=/tmp/model.ireeprof\n"
      "```\n"
      "\n"
      "Show raw counter samples for one dispatch event:\n"
      "\n"
      "```bash\n"
      "iree-profile counter --format=jsonl --counter_samples \\\n"
      "  --id=<dispatch_event_id> /tmp/model.ireeprof | \\\n"
      "  jq 'select(.type==\"counter_sample\") | \\\n"
      "      {sample_id,dispatch_event_id,key,counter,value,values}'\n"
      "```\n"
      "\n"
      "Inspect allocation high-water by physical device:\n"
      "\n"
      "```bash\n"
      "iree-profile memory --format=jsonl /tmp/model.ireeprof | \\\n"
      "  jq 'select(.type==\"memory_device\") | \\\n"
      "      {physical_device_ordinal,queue_inflight_current_bytes,"
      "queue_inflight_high_water_bytes,pool_reserved_current_bytes,"
      "pool_materialized_high_water_bytes}'\n"
      "```\n"
      "\n"
      "Find pools/providers with the largest memory high-water:\n"
      "\n"
      "```bash\n"
      "iree-profile memory --format=jsonl /tmp/model.ireeprof | \\\n"
      "  jq -s 'map(select(.type==\"memory_pool\")) | \\\n"
      "         sort_by(.lifecycle_high_water_bytes) | reverse[:20] | \\\n"
      "         map({kind,physical_device_ordinal,pool_id,memory_type,"
      "lifecycle_high_water_bytes,waits})'\n"
      "```\n"
      "\n"
      "Trace one allocation lifecycle:\n"
      "\n"
      "```bash\n"
      "iree-profile memory --format=jsonl --id=<allocation_id> \\\n"
      "  /tmp/model.ireeprof | \\\n"
      "  jq 'select(.type==\"memory_event\" or "
      ".type==\"memory_allocation\")'\n"
      "```\n"
      "\n"
      "Map executable ids and function ordinals to names:\n"
      "\n"
      "```bash\n"
      "iree-profile executable --format=jsonl /tmp/model.ireeprof | \\\n"
      "  jq 'select(.type==\"executable_function\") | \\\n"
      "      "
      "{executable_id,function_ordinal,key,binding_count,workgroup_size}'\n"
      "```\n",
      file);
}

static void iree_profile_print_agent_markdown_notes(FILE* file) {
  fputs(
      "\n"
      "## Notes For Agents\n"
      "\n"
      "Treat `summary` as the first sanity check: if clock fitting is "
      "unavailable\n"
      "or dispatch timestamps are invalid, timing conclusions are suspect. "
      "For\n"
      "automated analysis, prefer entering through `queue`, `command`,\n"
      "`executable`, or `dispatch` with JSONL and then drilling into\n"
      "`--dispatch_events` using the ids above. Do not infer queue ordering "
      "from\n"
      "record order alone; HAL queue ordering is established by semaphore "
      "edges,\n"
      "and full queue-edge records require producer support beyond the "
      "current\n"
      "dispatch-derived projections.\n",
      file);
}

void iree_profile_print_agent_markdown(FILE* file) {
  iree_profile_print_agent_markdown_capture(file);
  iree_profile_print_agent_markdown_commands(file);
  iree_profile_print_agent_markdown_jsonl_contracts(file);
  iree_profile_print_agent_markdown_jsonl_record_types(file);
  iree_profile_print_agent_markdown_cross_references(file);
  iree_profile_print_agent_markdown_recipes(file);
  iree_profile_print_agent_markdown_notes(file);
}

const char* iree_profile_usage_text(void) { return kIreeProfileUsage; }
