Alex-Vasile · Nov 6, 2026 · Nov 6, 2026
@@ -0,0 +1,97 @@
+#!/usr/bin/env python3
+# Copyright Advanced Micro Devices, Inc., or its affiliates.
+# SPDX-License-Identifier: MIT
+"""Compare two per-config timing passes side by side.
+
+Each argument is the OUT directory of a run_per_config.sh or run_per_config_build.sh pass: it holds
+AGGREGATE_analysis.txt (written by scripts/aggregate_timing.py) and one <config>/times.txt per
+config. The script prints the "TIMING BY PURPOSE" table of both aggregates side by side with the
+change in percent, then the summed pytest wall time, the slowest configs and every config that
+got more than 2 s slower.
+
+Usage: compare_aggregates.py BEFORE_DIR AFTER_DIR [--max-depth N]
+"""
+
+import argparse
+import glob
+import os
+import re
+
+
+def purpose_rows(path):
+    """(indent, category, milliseconds) for each row of the TIMING BY PURPOSE table."""
+    rows = []
+    inside = False
+    with open(path) as f:
+        for line in f:
+            if line.startswith("TIMING BY PURPOSE"):
+                inside = True
+                continue
+            if inside and line.startswith("TIMING BY TIMER NESTING"):
+                break
+            m = re.match(r"^(\s+)(\S.*?)\s{2,}(?:([\d,]+)\s+)?([\d,]+\.\d\d)\s", line)
+            if inside and m:
+                rows.append((len(m.group(1)), m.group(2).strip(), float(m.group(4).replace(",", ""))))
+    return rows
+
+
+def keyed(rows):
+    """Rows keyed by their path in the category tree, so equal names at different places differ."""
+    out = {}
+    stack = []
+    for indent, name, ms in rows:
+        while stack and stack[-1][0] >= indent:
+            stack.pop()
+        stack.append((indent, name))
+        out["/".join(n for _, n in stack)] = (indent, ms)
+    return out
+
+
+def pytest_seconds(directory):
+    """Pytest wall seconds per config, from <config>/times.txt."""
+    seconds = {}
+    for path in glob.glob(os.path.join(directory, "*", "times.txt")):
+        with open(path) as f:
+            m = re.search(r"seconds=(\d+)", f.read())
+        if m:
+            seconds[os.path.basename(os.path.dirname(path))] = int(m.group(1))
+    return seconds
+
+
+def main():
+    parser = argparse.ArgumentParser(description=__doc__, formatter_class=argparse.RawDescriptionHelpFormatter)
+    parser.add_argument("before", help="OUT directory of the 'before' pass")
+    parser.add_argument("after", help="OUT directory of the 'after' pass")
+    parser.add_argument("--max-depth", type=int, default=8, help="deepest category level printed")
+    args = parser.parse_args()
+
+    a = keyed(purpose_rows(os.path.join(args.before, "AGGREGATE_analysis.txt")))
+    b = keyed(purpose_rows(os.path.join(args.after, "AGGREGATE_analysis.txt")))
+    print(f"{'Category':70s} {'before s':>10s} {'after s':>10s} {'delta':>8s}")
+    for key in list(a) + [k for k in b if k not in a]:
+        indent = (a.get(key) or b.get(key))[0]
+        if (indent - 4) // 2 > args.max_depth:
+            continue
+        x = a.get(key, (0, 0))[1] / 1000
+        y = b.get(key, (0, 0))[1] / 1000
+        if max(x, y) < 1:
+            continue
+        delta = f"{(y - x) / x * 100:+.1f}%" if x else "new"
+        width = 70 - indent + 4
+        print(f"{' ' * (indent - 4)}{key.split('/')[-1][:width]:{width}s} {x:10.1f} {y:10.1f} {delta:>8s}")
+
+    tb = pytest_seconds(args.before)
+    ta = pytest_seconds(args.after)
+    total_b = sum(tb.values())
+    total_a = sum(ta.values())
+    change = f" ({(total_a - total_b) / total_b * 100:+.1f}%)" if total_b else ""
+    print(f"\npytest wall sum: before {total_b} s, after {total_a} s{change}")
+    print("largest per-config (before -> after s):")
+    for k in sorted(tb, key=lambda k: -tb[k])[:12]:
+        print(f"  {k:60s} {tb[k]:5d} -> {ta.get(k, 0):5d}")
+    slower = sorted(((k, tb[k], ta[k]) for k in tb if k in ta and ta[k] > tb[k] + 2), key=lambda t: t[1] - t[2])
+    print("slower by more than 2 s:", slower if slower else "none")
+
+
+if __name__ == "__main__":
+    main()
@@ -0,0 +1,83 @@
+#!/bin/bash
+# Copyright Advanced Micro Devices, Inc., or its affiliates.
+# SPDX-License-Identifier: MIT
+#
+# One side ("before" or "after") of a before/after timing of the common tests. Start one instance
+# per side, at the same time, each with its own source tree, client, cores and GPU, so that both
+# sides see the same host load. Each side runs:
+#   1. split/run_split.sh: the CI-like build phase then run phase (uninstrumented, -n 4), and
+#      records the configs whose run passed (<side>/split_run_passed.txt);
+#   2. waits until every side in SIDES has finished step 1, then writes the shared config list
+#      (<side>/configs.txt): the union of the configs that passed in any side's split run phase;
+#   3. per_config/run_per_config_build.sh over that list, instrumented, caches cleared once before
+#      the first config, artifacts kept;
+#   4. per_config/run_per_config.sh on those artifacts, instrumented, caches not cleared.
+# Compare the two sides with before_after/compare_aggregates.py.
+# Usage: bash run_side.sh <side>        (<side> is one of SIDES)
+# Env:   ROOT        work directory shared by the sides (required); <side>'s output goes to
+#                    $ROOT/<side>/
+#        TL, CLIENT  this side's TensileLite tree and prebuilt tensilelite-client (required)
+#        SIDES       side names, default "before after"
+#        GPU_TARGET, CORES, GPU, ANALYZE_DIR, DIST, CLEAR_RUN_CACHE: see ../measure_env.sh and
+#                    ../split/run_split.sh. Give each side its own CORES and GPU.
+set -u
+S=${1:?usage: run_side.sh <side>}
+: "${ROOT:?set ROOT}" "${TL:?set TL}" "${CLIENT:?set CLIENT}"
+SIDES=${SIDES:-before after}
+case " $SIDES " in
+  *" $S "*) ;;
+  *) echo "side '$S' is not in SIDES='$SIDES'" >&2; exit 1 ;;
+esac
+SCRIPTS=$(dirname "$(dirname "$(realpath -- "${BASH_SOURCE[0]}")")")
+export TL CLIENT
+R=$ROOT/$S
+mkdir -p "$R"
+export TMPDIR=$R/tmp
+mkdir -p "$TMPDIR"
+log() { echo "$(date -u +%FT%TZ) $*" >> "$R/driver.log"; }
+# Host conditions at each step, for judging how comparable the two sides are on a shared host.
+snap() {
+  {
+    echo "== $1 $(date -u +%FT%TZ)"
+    uptime
+    rocm-smi -d "${GPU:-0}" --showuse --showmemuse --showpids 2>/dev/null | grep -Ev '^=+$|^$'
+  } >> "$R/conditions.log"
+}
+
+# 1. CI-like split: wall time of each phase and the configs that pass
+if ! grep -qs ALLDONE "$R/split/times.txt"; then
+  snap split_start; log split start
+  OUT=$R/split ART=$R/split_art TENSILE_CACHE=$R/cache_split/tensile \
+    TENSILELITE_CACHE=$R/cache_split/tensilelite bash "$SCRIPTS/split/run_split.sh"
+  snap split_end; log split done
+fi
+python - "$R/split/run.xml" > "$R/split_run_passed.txt" <<'EOF'
+import re
+import sys
+import xml.etree.ElementTree as ET
+
+for tc in ET.parse(sys.argv[1]).iter("testcase"):
+    m = re.search(r"test_config_run\[(.*)\]", tc.get("name", ""))
+    if m and not any(c.tag in ("failure", "error", "skipped") for c in tc):
+        print(m.group(1))
+EOF
+touch "$R/split.done"
+
+# 2. shared config list: union of the configs whose run passed in any side's split run phase
+for s in $SIDES; do
+  while [ ! -f "$ROOT/$s/split.done" ]; do sleep 30; done
+done
+for s in $SIDES; do cat "$ROOT/$s/split_run_passed.txt"; done | sort -u > "$R/configs.txt"
+log "configs: $(wc -l < "$R/configs.txt")"
+
+# 3. per-config build, instrumented; caches cleared once at the start, artifacts kept
+export ART=$R/art TENSILE_CACHE=$R/cache/tensile TENSILELITE_CACHE=$R/cache/tensilelite
+snap build_start; log build start
+OUT=$R/build SCRATCH=$R/scratch CLEAR_CACHE=once KEEP_ARTIFACTS=1 \
+  bash "$SCRIPTS/per_config/run_per_config_build.sh" "$R/configs.txt"
+snap build_end; log build done
+
+# 4. per-config run on those artifacts, caches not cleared
+snap run_start; log run start
+OUT=$R/run CLEAR_CACHE=never bash "$SCRIPTS/per_config/run_per_config.sh" "$R/configs.txt"
+snap run_end; log run done
@@ -0,0 +1,93 @@
+# Copyright Advanced Micro Devices, Inc., or its affiliates.
+# SPDX-License-Identifier: MIT
+#
+# Sourced by split/run_split.sh, per_config/run_per_config.sh and per_config/run_per_config_build.sh.
+# Every location the measurement scripts read, write or delete can be overridden from the
+# environment; the defaults are the paths the committed baselines were measured with.
+#   TL           TensileLite source tree pytest runs from (e.g. a git worktree of another commit)
+#   CLIENT       prebuilt tensilelite-client
+#   ART          build artifact dir (split build phase writes it, the run phases read it). The
+#                default is the measured baseline, which check_rm_target refuses to delete, so
+#                split/run_split.sh needs ART set to a new dir.
+#   TENSILE_CACHE      Tensile's cache root, cleared by clear_caches. TENSILE_HELPER_CACHE_DIR,
+#                      TENSILE_HELPER_PCH_DIR and TENSILE_HELPER_CHUNK_CACHE_DIR are always
+#                      exported as $TENSILE_CACHE/helper_cache, $TENSILE_CACHE/helper_pch and
+#                      $TENSILE_CACHE/helper_chunk_cache (inherited values are ignored), so clearing
+#                      TENSILE_CACHE clears the helper-kernel cache, HIP-header PCH and helper chunk
+#                      cache Tensile really uses. TENSILE_DISABLE_HELPER_PCH=1 (compiles without a
+#                      PCH) and TENSILE_DISABLE_HELPER_CHUNK_CACHE=1 (compiles every helper chunk)
+#                      are passed through.
+#   TENSILELITE_CACHE  ~/.tensilelite, cleared by clear_caches as before (current TensileLite
+#                      sources neither read nor write it).
+#   GPU_TARGET   --gpu-targets of the build phases (default gfx950)
+#   CORES        taskset core list (default 0-24; OMP_NUM_THREADS stays 25, so give 25 cores)
+#   GPU          ROCR_VISIBLE_DEVICES of the runs (default 0)
+#   ANALYZE_DIR  directory holding analyze_timing.py / aggregate_timing.py (default $TL/scripts),
+#                so logs of an older TL can be analysed with a newer analyzer
+#   CLEAR_CACHE  per-config | once | never: when the per-config scripts clear the caches (default
+#                per-config; KEEP_CACHE=1 means once). once clears before the first config of the
+#                pass; never leaves the caches as the previous pass left them.
+# OUT, LIST and SCRATCH are set by the scripts themselves (their defaults depend on the script);
+# each script then calls check_rm_target on every directory it deletes.
+TL=${TL:-/home/alvasile/rocm-libraries/projects/hipblaslt/tensilelite}
+CLIENT=${CLIENT:-/tmp/tl_client_msgpack/tensilelite/client/tensilelite-client}
+ART=${ART:-/tmp/tl_split_gfx950/artifacts}
+TENSILE_CACHE=${TENSILE_CACHE:-/home/alvasile/.tensile}
+TENSILELITE_CACHE=${TENSILELITE_CACHE:-/home/alvasile/.tensilelite}
+GPU_TARGET=${GPU_TARGET:-gfx950}
+CORES=${CORES:-0-24}
+GPU=${GPU:-0}
+ANALYZE_DIR=${ANALYZE_DIR:-$TL/scripts}
+if [ -z "$CLEAR_CACHE" ]; then
+  if [ "${KEEP_CACHE:-0}" = "1" ]; then CLEAR_CACHE=once; else CLEAR_CACHE=per-config; fi
+fi
+case $CLEAR_CACHE in
+  per-config|once|never) ;;
+  *) echo "CLEAR_CACHE='$CLEAR_CACHE' must be per-config, once or never" >&2; exit 1 ;;
+esac
+export TENSILE_HELPER_CACHE_DIR=$TENSILE_CACHE/helper_cache
+export TENSILE_HELPER_PCH_DIR=$TENSILE_CACHE/helper_pch
+export TENSILE_HELPER_CHUNK_CACHE_DIR=$TENSILE_CACHE/helper_chunk_cache
+MEASURE_DIR=$(dirname "$(realpath -- "${BASH_SOURCE[0]}")")
+# Measured baseline client build and artifacts; never deleted.
+BASELINE_DIRS=(/tmp/tl_client_msgpack /tmp/tl_split_gfx950/artifacts)
+
+# check_rm_target VAR...: exit unless every $VAR is safe to `rm -rf`: an absolute path at least two
+# levels deep that is not $HOME or a directory containing it, and that neither is, contains, nor is
+# inside $TL, $OUT, this script's directory, a baseline dir, $CLIENT's directory or $ART (unless
+# VAR is ART).
+check_rm_target() {
+  local var val real keep keeps
+  for var in "$@"; do
+    val=${!var}
+    case $val in
+      /*) ;;
+      *) echo "$var='$val' must be an absolute path" >&2; exit 1 ;;
+    esac
+    real=$(realpath -m -- "$val")
+    case $real in
+      /*/*) ;;
+      *) echo "$var='$val' is too broad to delete" >&2; exit 1 ;;
+    esac
+    if [ -n "$HOME" ]; then
+      keep=$(realpath -m -- "$HOME")
+      case $keep/ in
+        "$real"/*) echo "$var='$val' is or contains $keep; refusing to delete it" >&2; exit 1 ;;
+      esac
+    fi
+    keeps=("$TL" "$OUT" "$MEASURE_DIR" "${BASELINE_DIRS[@]}" "$(dirname -- "$CLIENT")")
+    [ "$var" = ART ] || keeps+=("$ART")
+    for keep in "${keeps[@]}"; do
+      [ -n "$keep" ] || continue
+      keep=$(realpath -m -- "$keep")
+      case $keep/ in
+        "$real"/*) echo "$var='$val' is or contains $keep; refusing to delete it" >&2; exit 1 ;;
+      esac
+      case $real/ in
+        "$keep"/*) echo "$var='$val' is inside $keep; refusing to delete it" >&2; exit 1 ;;
+      esac
+    done
+  done
+}
+
+clear_caches() { rm -rf -- "$TENSILE_CACHE" "$TENSILELITE_CACHE"; }
@@ -0,0 +1,64 @@
+#!/bin/bash
+# Copyright Advanced Micro Devices, Inc., or its affiliates.
+# SPDX-License-Identifier: MIT
+#
+# Run step of every gfx950 common config that passed in the split run phase, one at a time
+# (serial, no xdist), with TimingInstrumentation on, and save per config:
+#   <name>/<name>.log.gz            raw instrumented pytest/client log
+#   <name>/<name>_analysis.txt      scripts/analyze_timing.py output
+#   <name>/times.txt                pytest wall seconds + return code
+# Build artifacts come from the split build phase (../split) or from
+# run_per_config_build.sh with KEEP_ARTIFACTS=1. Caches (~/.tensile, ~/.tensilelite) are cleared
+# before each config, as in run_cold_warm.sh / the split phases, unless CLEAR_CACHE=once
+# (or KEEP_CACHE=1): then they are cleared once before the first config and kept between configs
+# (a normal run, where e.g. the helper-kernel cache is shared by all configs). CLEAR_CACHE=never
+# keeps the cache a preceding build pass left (a CI-like build-then-run on one node).
+# When every config is DONE, scripts/aggregate_timing.py sums the logs of ALL configs (failed ones
+# included) into $OUT/AGGREGATE_analysis.txt. Re-running the script after a finished pass skips the
+# runs and only regenerates the aggregate.
+# Usage: bash run_per_config.sh [configs.txt]   (one tensilelite/Tests/... yaml path per line)
+# Env: TL, CLIENT, ART, TENSILE_CACHE, TENSILELITE_CACHE (see ../measure_env.sh),
+#      CORES, GPU, ANALYZE_DIR, CLEAR_CACHE (see ../measure_env.sh),
+#      OUT (default $TL/timing_runs/full_runs_gfx950/per_config), KEEP_CACHE.
+source "$(dirname "${BASH_SOURCE[0]}")/../measure_env.sh" || exit 1
+OUT=${OUT:-$TL/timing_runs/full_runs_gfx950/per_config}
+LIST=${1:-$OUT/configs.txt}
+check_rm_target TENSILE_CACHE TENSILELITE_CACHE
+cd "$TL" || exit 1
+mkdir -p "$OUT"
+export OMP_NUM_THREADS=25 OPENBLAS_NUM_THREADS=25 ROCR_VISIBLE_DEVICES=$GPU PYTHONUNBUFFERED=1
+TO="--cxx-compiler,/usr/bin/amdclang++,--global-parameters,TimingInstrumentation=True"
+COMMON=(-s -v -n 0 --timeout=5400 -m "common and not skip-gfx94X"
+  "--global-parameters=LibraryFormat='msgpack'" "--tensile-options=$TO" -p no:warnings -p no:cacheprovider)
+
+n=0; total=$(wc -l < "$LIST")
+if [ "$CLEAR_CACHE" = once ]; then
+  clear_caches
+  echo "CLEAR_CACHE=once: caches cleared once at $(date -u +%FT%TZ)" >> "$OUT/progress.txt"
+elif [ "$CLEAR_CACHE" = never ]; then
+  echo "CLEAR_CACHE=never: caches kept at $(date -u +%FT%TZ)" >> "$OUT/progress.txt"
+fi
+while read -r cfg; do
+  [ -z "$cfg" ] && continue
+  n=$((n+1))
+  name=$(echo "${cfg#tensilelite/Tests/}" | sed -e 's/\.yaml$//' -e 's#/#__#g')
+  D=$OUT/$name
+  if [ -f "$D/times.txt" ] && grep -q DONE "$D/times.txt"; then continue; fi   # resumable
+  mkdir -p "$D"
+  RID="tensilelite/Tests/common/test_config_run.py::test_config_run[$cfg]"
+  if [ "$CLEAR_CACHE" = per-config ]; then
+    clear_caches
+  fi
+  s=$(date +%s)
+  taskset -c "$CORES" pytest "$RID" --use-cache --artifact-dir "$ART" --prebuilt-client="$CLIENT" \
+    "${COMMON[@]}" > "$D/$name.log" 2>&1
+  rc=$?
+  echo "run rc=$rc seconds=$(( $(date +%s)-s )) ($n/$total)" > "$D/times.txt"
+  python "$ANALYZE_DIR/analyze_timing.py" "$D/$name.log" > "$D/${name}_analysis.txt" 2>&1
+  gzip -f "$D/$name.log"
+  echo DONE >> "$D/times.txt"
+  echo "$(date -u +%T) $n/$total $cfg rc=$rc" >> "$OUT/progress.txt"
+done < "$LIST"
+python "$ANALYZE_DIR/aggregate_timing.py" "$OUT"/*/*.log.gz > "$OUT/AGGREGATE_analysis.txt" 2>&1
+echo "aggregate rc=$? logs=$(ls "$OUT"/*/*.log.gz | wc -l)" >> "$OUT/progress.txt"
+echo ALLDONE > "$OUT/done.txt"