lablab-ai-amd-developer-hackathon/gpu-goblin

1"""profile_parser — turn goblin_runner.sh artefacts into RunMetrics.2 3Reads three files written by goblin_runner.sh (architecture.md §5):4 5    <out_dir>/trace.csv          rocprofv3 kernel trace6    <out_dir>/torch_profile.json torch.profiler chrome trace7    <out_dir>/amd_smi.csv        amd-smi telemetry, ~200 ms cadence8 9and produces one `RunMetrics` with a populated `WasteBudget`.10 11Each waste-budget bucket is computed with a deliberately simple, documented12heuristic — these are best-effort signals for the agent, NOT measured13ground-truth. See the docstring on `_waste_budget()` for the per-bucket14formula and its known failure modes.15 16This module is import-tolerant on machines without the rocprofv3 stack —17it only reads files. Missing or unparseable files degrade individual18metrics to zero and append a warning to RunMetrics.warnings rather than19raising. LiveRunner ultimately decides what to do with parse failures.20"""21 22from __future__ import annotations23 24import csv25import json26import logging27import re28from dataclasses import dataclass29from pathlib import Path30from typing import Any31 32from agent.schemas import KernelEntry, RunMetrics, WasteBudget, WorkloadConfig33 34_LOG = logging.getLogger(__name__)35 36 37# MI300X has 192 GB HBM3. We treat sustained <70% utilisation as headroom.38_HBM_TOTAL_GB = 192.039_HBM_HEALTHY_TARGET = 0.7040 41# Kernel-name patterns used by the heuristics.42#43# We use `(?<![A-Za-z0-9])` / `(?![A-Za-z0-9])` instead of `\b` so that44# underscores act as token separators — kernel names like45# `rccl_AllReduce` and `hipBLASLt_generic_gemm` should match.46_BOUND_L = r"(?<![A-Za-z0-9])"47_BOUND_R = r"(?![A-Za-z0-9])"48_RCCL_PATTERN = re.compile(49    _BOUND_L + r"(rccl|nccl|all[_-]?reduce|broadcast|reduce[_-]?scatter)" + _BOUND_R, re.I50)51_GEMM_PATTERN = re.compile(_BOUND_L + r"(gemm|matmul|hgemm|sgemm|hipblaslt)" + _BOUND_R, re.I)52_GENERIC_GEMM_PATTERN = re.compile(53    _BOUND_L + r"(generic|fallback|naive|reference)" + _BOUND_R, re.I54)55_FP16_PATTERN = re.compile(_BOUND_L + r"(fp16|half|f16)" + _BOUND_R, re.I)56_BF16_PATTERN = re.compile(_BOUND_L + r"(bf16|bfloat16)" + _BOUND_R, re.I)57 58 59# ---------------------------------------------------------------------------60# Public entry point61# ---------------------------------------------------------------------------62 63 64def parse(65    out_dir: Path,66    config: WorkloadConfig | None = None,67    steps: int = 10,68) -> RunMetrics:69    """Build a `RunMetrics` from goblin_runner.sh artefacts in `out_dir`.70 71    `config` is used by some heuristics (e.g. precision_path skips the72    estimate when the config is already bf16). It can be None at the cost73    of a more conservative waste-budget estimate.74 75    Always returns a RunMetrics object — individual metrics degrade to76    zero with a warning rather than raising, because LiveRunner ultimately77    decides whether parse failures should trigger fallback.78    """79    warnings: list[str] = []80 81    kernels = _read_kernels(out_dir / "trace.csv", warnings)82    torch_summary = _read_torch_profile(out_dir / "torch_profile.json", warnings)83    smi = _read_amd_smi(out_dir / "amd_smi.csv", warnings)84 85    top_kernels = _top_kernels(kernels, top_n=5)86    gpu_util_pct = smi.gpu_util_pct if smi.gpu_util_pct is not None else _gpu_util_from_kernels(87        kernels, torch_summary88    )89    # Write the resolved gpu_util back into smi so _waste_budget sees the90    # same number RunMetrics reports. Without this, a build where amd-smi's91    # `usage` column returns N/A leaves smi.gpu_util_pct=None — and92    # _waste_budget reads it as 0%, which forces host_gap to consume the93    # full step time and collapses kernel_shape / precision_path to zero94    # (both multiplied by gpu_util).95    smi.gpu_util_pct = gpu_util_pct96 97    waste_budget = _waste_budget(98        kernels=kernels,99        torch_summary=torch_summary,100        smi=smi,101        config=config,102    )103 104    attention_kernel_loaded = _detect_attention_kernel(kernels)105 106    return RunMetrics(107        steps=steps,108        tokens_per_sec=torch_summary.tokens_per_sec or 0.0,109        mfu_pct=torch_summary.mfu_pct or 0.0,110        hbm_peak_gb=smi.hbm_peak_gb or 0.0,111        hbm_avg_gb=smi.hbm_avg_gb or 0.0,112        gpu_util_pct=gpu_util_pct,113        top_kernels=top_kernels,114        attention_kernel_loaded=attention_kernel_loaded,115        waste_budget=waste_budget,116        warnings=warnings,117        rocm_version=smi.rocm_version or "unknown",118        pytorch_version=torch_summary.pytorch_version or "unknown",119        runner_kind="live",120    )121 122 123# ---------------------------------------------------------------------------124# Internal data carriers125# ---------------------------------------------------------------------------126 127 128@dataclass129class _Kernel:130    name: str131    duration_ns: int132    """Kernel duration in nanoseconds."""133 134    is_collective: bool = False135    is_gemm: bool = False136    is_generic_gemm: bool = False137 138 139@dataclass140class _TorchSummary:141    tokens_per_sec: float | None = None142    mfu_pct: float | None = None143    pytorch_version: str | None = None144    step_time_seconds: float | None = None145    """Average wall-clock seconds per training step (used by waste budget)."""146 147    host_busy_fraction: float | None = None148    """Fraction of step time the host (CPU) spent doing non-launch work.149    Heuristic: `cpu_self_time / step_time` reported by torch.profiler."""150 151 152@dataclass153class _SmiSummary:154    hbm_peak_gb: float | None = None155    hbm_avg_gb: float | None = None156    gpu_util_pct: float | None = None157    rocm_version: str | None = None158 159 160# ---------------------------------------------------------------------------161# trace.csv (rocprofv3) → list[_Kernel]162# ---------------------------------------------------------------------------163 164 165def _read_kernels(path: Path, warnings: list[str]) -> list[_Kernel]:166    """Parse rocprofv3 kernel-trace CSV into a list of `_Kernel` records.167 168    rocprofv3 column names vary slightly by version. We look up by header and169    accept the common aliases; if the file is missing or unparseable we170    return an empty list and append a warning so the caller can decide.171    """172    if not path.exists():173        warnings.append(f"profile_parser: kernel trace not found at {path}")174        return []175 176    try:177        with path.open(newline="") as f:178            reader = csv.DictReader(f)179            if reader.fieldnames is None:180                warnings.append(f"profile_parser: empty kernel trace at {path}")181                return []182            name_col = _pick_column(reader.fieldnames, ["KernelName", "Kernel_Name", "kernel_name", "Name"])183            start_col = _pick_column(reader.fieldnames, ["BeginNs", "start_ns", "BeginNS", "Start"])184            end_col = _pick_column(reader.fieldnames, ["EndNs", "end_ns", "EndNS", "End"])185            duration_col = _pick_column(reader.fieldnames, ["DurationNs", "duration_ns", "Duration"])186 187            kernels: list[_Kernel] = []188            for row in reader:189                name = (row.get(name_col) or "").strip() if name_col else ""190                if not name:191                    continue192                duration = _row_duration_ns(row, duration_col, start_col, end_col)193                if duration <= 0:194                    continue195                kernels.append(196                    _Kernel(197                        name=name,198                        duration_ns=duration,199                        is_collective=bool(_RCCL_PATTERN.search(name)),200                        is_gemm=bool(_GEMM_PATTERN.search(name)),201                        is_generic_gemm=bool(202                            _GEMM_PATTERN.search(name) and _GENERIC_GEMM_PATTERN.search(name)203                        ),204                    )205                )206            return kernels207    except (OSError, csv.Error) as exc:208        warnings.append(f"profile_parser: failed to read kernel trace ({exc})")209        return []210 211 212def _pick_column(fieldnames: list[str], candidates: list[str]) -> str | None:213    """Pick the first matching column name, with three fallback tiers:214        1. Exact match.215        2. Case-insensitive exact match.216        3. Substring match (case-insensitive) — every token of any candidate217           must appear somewhere in the field name. Tolerates the column-name218           drift between rocprofv3 / amd-smi versions (e.g. `VRAM_USED` vs219           `vram_used_mb` vs `VRAM USED MB`).220    """221    for c in candidates:222        if c in fieldnames:223            return c224    lower = {f.lower(): f for f in fieldnames}225    for c in candidates:226        if c.lower() in lower:227            return lower[c.lower()]228    # Substring tier: split each candidate on _/space, require all tokens229    # appear in the (lowercased) field name. Avoids matching too eagerly230    # by requiring every token of the candidate.231    for c in candidates:232        tokens = [t for t in c.lower().replace("_", " ").split() if t]233        if not tokens:234            continue235        for fname in fieldnames:236            fl = fname.lower()237            if all(t in fl for t in tokens):238                return fname239    return None240 241 242def _row_duration_ns(243    row: dict[str, str], duration_col: str | None, start_col: str | None, end_col: str | None244) -> int:245    if duration_col and row.get(duration_col):246        try:247            return int(float(row[duration_col]))248        except ValueError:249            return 0250    if start_col and end_col and row.get(start_col) and row.get(end_col):251        try:252            return int(float(row[end_col])) - int(float(row[start_col]))253        except ValueError:254            return 0255    return 0256 257 258def _top_kernels(kernels: list[_Kernel], top_n: int) -> list[KernelEntry]:259    if not kernels:260        return []261    total = sum(k.duration_ns for k in kernels) or 1262    by_name: dict[str, int] = {}263    for k in kernels:264        by_name[k.name] = by_name.get(k.name, 0) + k.duration_ns265    ranked = sorted(by_name.items(), key=lambda kv: kv[1], reverse=True)[:top_n]266    return [KernelEntry(name=name, pct_time=ns / total * 100.0) for name, ns in ranked]267 268 269def _detect_attention_kernel(kernels: list[_Kernel]) -> str:270    for k in kernels:271        n = k.name.lower()272        if "flash_attn" in n and "rocm" in n:273            return "flash_rocm"274        if "flash" in n and "attn" in n:275            return "flash"276        if "scaled_dot_product_attention" in n:277            return "sdpa"278    if kernels:279        return "eager"  # nothing flash-shaped; default conservative label280    return "unknown"281 282 283def _gpu_util_from_kernels(kernels: list[_Kernel], torch_summary: _TorchSummary) -> float:284    """Fallback GPU util when amd-smi is missing.285 286    util ≈ sum(kernel duration) / total wall-clock step time.287    """288    if not kernels or torch_summary.step_time_seconds in (None, 0):289        return 0.0290    total_kernel_ns = sum(k.duration_ns for k in kernels)291    wall_ns = torch_summary.step_time_seconds * 1e9292    if wall_ns <= 0:293        return 0.0294    return min(100.0, total_kernel_ns / wall_ns * 100.0)295 296 297# ---------------------------------------------------------------------------298# torch_profile.json (torch.profiler chrome trace) → _TorchSummary299# ---------------------------------------------------------------------------300 301 302def _read_torch_profile(path: Path, warnings: list[str]) -> _TorchSummary:303    """Pull tokens/sec, MFU, and step timing from a torch.profiler artefact.304 305    The user script (workloads/train_qwen_lora.py in Phase 3) is306    responsible for embedding `tokens_per_sec`, `mfu_pct`, `pytorch_version`307    and `step_time_seconds` in the trace as `metadata` events. If those are308    missing, we estimate `step_time_seconds` from the total trace duration.309    """310    summary = _TorchSummary()311    if not path.exists():312        warnings.append(f"profile_parser: torch profile not found at {path}")313        return summary314    try:315        data = json.loads(path.read_text())316    except (OSError, json.JSONDecodeError) as exc:317        warnings.append(f"profile_parser: failed to read torch profile ({exc})")318        return summary319 320    # torch.profiler chrome trace is `{"traceEvents": [...], "metadata": {...}}`321    metadata = data.get("metadata") if isinstance(data, dict) else None322    if isinstance(metadata, dict):323        summary.tokens_per_sec = _coerce_float(metadata.get("tokens_per_sec"))324        summary.mfu_pct = _coerce_float(metadata.get("mfu_pct"))325        summary.pytorch_version = metadata.get("pytorch_version") or metadata.get("torch_version")326        summary.step_time_seconds = _coerce_float(metadata.get("step_time_seconds"))327        summary.host_busy_fraction = _coerce_float(metadata.get("host_busy_fraction"))328 329    events = data.get("traceEvents") if isinstance(data, dict) else None330    if isinstance(events, list):331        if summary.step_time_seconds is None:332            summary.step_time_seconds = _step_time_from_events(events)333        if summary.host_busy_fraction is None:334            summary.host_busy_fraction = _host_busy_from_events(events)335 336    return summary337 338 339def _coerce_float(v: Any) -> float | None:340    if v is None:341        return None342    try:343        return float(v)344    except (TypeError, ValueError):345        return None346 347 348def _step_time_from_events(events: list[dict]) -> float | None:349    """Estimate per-step wall-clock seconds from chrome-trace duration events.350 351    Looks for `name == "ProfilerStep#*"` complete events; falls back to the352    overall trace span if those aren't present.353    """354    durations: list[float] = []355    overall_start: float | None = None356    overall_end: float | None = None357    for ev in events:358        if not isinstance(ev, dict):359            continue360        name = ev.get("name", "")361        ts = ev.get("ts")362        dur = ev.get("dur")363        if isinstance(name, str) and name.startswith("ProfilerStep") and isinstance(dur, (int, float)):364            durations.append(float(dur) / 1e6)  # us → seconds365        if isinstance(ts, (int, float)) and isinstance(dur, (int, float)):366            start = float(ts)367            end = start + float(dur)368            overall_start = start if overall_start is None else min(overall_start, start)369            overall_end = end if overall_end is None else max(overall_end, end)370    if durations:371        return sum(durations) / len(durations)372    if overall_start is not None and overall_end is not None and overall_end > overall_start:373        return (overall_end - overall_start) / 1e6374    return None375 376 377def _host_busy_from_events(events: list[dict]) -> float | None:378    """Heuristic: cpu_op event time / total event span.379 380    Used by the data_wait waste bucket to disambiguate "GPU idle because the381    host is busy preparing the next batch" from "GPU idle because nothing382    is running anywhere".383    """384    cpu_op_us = 0.0385    span_min: float | None = None386    span_max: float | None = None387    for ev in events:388        if not isinstance(ev, dict):389            continue390        cat = ev.get("cat", "")391        ts = ev.get("ts")392        dur = ev.get("dur")393        if not (isinstance(ts, (int, float)) and isinstance(dur, (int, float))):394            continue395        ts_f = float(ts)396        dur_f = float(dur)397        end = ts_f + dur_f398        span_min = ts_f if span_min is None else min(span_min, ts_f)399        span_max = end if span_max is None else max(span_max, end)400        if isinstance(cat, str) and "cpu_op" in cat.lower():401            cpu_op_us += dur_f402    if span_min is None or span_max is None or span_max <= span_min:403        return None404    span = span_max - span_min405    if span <= 0:406        return None407    return min(1.0, cpu_op_us / span)408 409 410# ---------------------------------------------------------------------------411# amd_smi.csv → _SmiSummary412# ---------------------------------------------------------------------------413 414 415def _read_amd_smi(path: Path, warnings: list[str]) -> _SmiSummary:416    """Aggregate amd-smi polling output into HBM peak/avg + GPU util."""417    summary = _SmiSummary()418    if not path.exists():419        warnings.append(f"profile_parser: amd-smi telemetry not found at {path}")420        return summary421 422    try:423        raw = path.read_text()424    except OSError as exc:425        warnings.append(f"profile_parser: failed to read amd-smi telemetry ({exc})")426        return summary427 428    csv_text = _strip_amd_smi_preamble(raw)429    if csv_text is None:430        warnings.append(f"profile_parser: no parseable header in amd-smi telemetry at {path}")431        return summary432 433    try:434        import io as _io435 436        reader = csv.DictReader(_io.StringIO(csv_text))437        if reader.fieldnames is None:438            warnings.append(f"profile_parser: empty amd-smi telemetry at {path}")439            return summary440        hbm_col = _pick_column(441            reader.fieldnames,442            [443                # `amd-smi metric --mem-usage --csv` (ROCm 7.x) emits444                # "used_vram" in MB. Older `amd-smi monitor --vram-usage`445                # emits "VRAM_USED" / "vram_used_mb" depending on minor446                # version.447                "used_vram",448                "USED_VRAM",449                "VRAM_USED_MB",450                "vram_used_mb",451                "VRAM_USED",452                "VRAM_USED_GB",453                "vram_used",454                "VRAM Used",455                # Older rocm 6.x naming456                "MEM_USED",457                "mem_used",458            ],459        )460        util_col = _pick_column(461            reader.fieldnames,462            [463                # `amd-smi metric --usage --csv` (ROCm 7.x) emits a single464                # consolidated "usage" column. Some builds report N/A here465                # — _coerce_float drops those silently and we fall back to466                # the kernel-trace gpu_util estimate downstream.467                "usage",468                "USAGE",469                # `amd-smi monitor --gfx` (ROCm 7.x) → "gfx_util"470                "GFX_UTIL",471                "gfx_util",472                "GFX_UTILIZATION",473                "gfx_utilization",474                # rocm 6.x and older475                "GFX_ACTIVITY",476                "gfx_activity",477                "GPU_USE",478                "GFX %",479                "Util",480            ],481        )482        rocm_col = _pick_column(reader.fieldnames, ["ROCM_VERSION", "rocm_version"])483 484        hbm_samples: list[float] = []485        util_samples: list[float] = []486        for row in reader:487            if hbm_col:488                hbm_gb = _hbm_to_gb(row.get(hbm_col), hbm_col)489                if hbm_gb is not None:490                    hbm_samples.append(hbm_gb)491            if util_col:492                util = _coerce_float(row.get(util_col))493                if util is not None:494                    util_samples.append(min(100.0, util))495            if rocm_col and summary.rocm_version is None:496                val = (row.get(rocm_col) or "").strip()497                if val:498                    summary.rocm_version = val499        if hbm_samples:500            summary.hbm_peak_gb = max(hbm_samples)501            summary.hbm_avg_gb = sum(hbm_samples) / len(hbm_samples)502        if util_samples:503            summary.gpu_util_pct = sum(util_samples) / len(util_samples)504        return summary505    except csv.Error as exc:506        warnings.append(f"profile_parser: failed to parse amd-smi telemetry ({exc})")507        return summary508 509 510def _strip_amd_smi_preamble(raw: str) -> str | None:511    """Drop everything before the first real CSV header and dedupe repeated512    header lines — both noise produced by `amd-smi <subcmd> --watch`.513 514    --watch prints a "'CTRL' + 'C' to stop watching output:" banner once at515    the top, then re-emits the CSV header on every iteration. csv.DictReader516    naively reads the banner as fieldnames and treats every subsequent517    header as a misshapen data row. Pre-strip both before handing it off.518 519    Returns a CSV string ready for DictReader, or None if no header line520    is recognisable.521    """522    lines = raw.splitlines()523    header_idx: int | None = None524    for i, line in enumerate(lines):525        if "," not in line:526            continue527        lower = line.lower()528        # Recognised tokens come from the columns amd-smi metric / monitor529        # actually emit. Match conservatively — banners or other noise can530        # contain commas too.531        if any(tok in lower for tok in ("vram", "gfx_", "timestamp,", "gpu_use", "gpu,")):532            header_idx = i533            break534    if header_idx is None:535        return None536 537    header = lines[header_idx]538    data = [539        line540        for line in lines[header_idx + 1 :]541        if line.strip() and line != header542    ]543    return header + "\n" + "\n".join(data) + ("\n" if data else "")544 545 546def _hbm_to_gb(raw: str | None, column_name: str | None = None) -> float | None:547    """amd-smi sometimes reports VRAM in MB, sometimes in GB.548 549    First check the column name — `*vram*` / `*_mb` columns are MB-typed550    in every amd-smi build we've seen; `*_gb` is GB. Without a column-name551    hint, fall back to a value heuristic. The old "v > 1024 ⇒ MB" heuristic552    misclassified small idle samples (e.g. 285 MB at GPU idle) as GB and553    inflated the peak across the run, which then forced memory_headroom to554    zero downstream.555    """556    if not raw:557        return None558    try:559        v = float(str(raw).strip().replace("MB", "").replace("GB", "").replace(",", ""))560    except ValueError:561        return None562    if column_name:563        lower = column_name.lower()564        if "_mb" in lower or lower.endswith("mb"):565            return v / 1024.0566        if "_gb" in lower or lower.endswith("gb"):567            return v568        if "vram" in lower or "mem" in lower:569            # amd-smi metric / monitor default to MB for the bare570            # `used_vram` / `mem_used` columns on every ROCm 6.x+ build.571            return v / 1024.0572    # No column hint — fall back to value heuristic.573    if v > 1024.0:574        return v / 1024.0575    return v576 577 578# ---------------------------------------------------------------------------579# Waste-budget heuristics580# ---------------------------------------------------------------------------581 582 583def _waste_budget(584    *,585    kernels: list[_Kernel],586    torch_summary: _TorchSummary,587    smi: _SmiSummary,588    config: WorkloadConfig | None,589) -> WasteBudget:590    """Decompose step time into the seven WasteBudget buckets (architecture.md §3).591 592    These are HEURISTICS, not measurements. Each bucket is in seconds-per-step593    so they can be summed against `step_time_seconds`. If we can't observe a594    bucket we set it to 0 — `evidence_coverage` in `propose_patch` then595    discounts confidence accordingly.596 597    Per-bucket logic:598 599      data_wait600        Fraction of step time where GPU util was below 30% AND the host was601        busy (host_busy_fraction > 0.5). Maps to dataloader / H2D copy stalls.602 603      precision_path604        If the user is already on bf16/fp8 we skip this bucket (recovery is 0).605        Otherwise we estimate from kernel names: time spent in fp16-tagged606        GEMMs is the recoverable surface; bf16-tagged kernels are not.607 608      kernel_shape609        Fraction of total GEMM kernel time spent on kernels whose names610        match `generic|fallback|naive|reference` — i.e. hipBLASLt/MIOpen611        couldn't pick a tuned tile size and fell back to a slow path.612 613      host_gap614        Time the GPU was idle while the host was NOT busy either — pure615        launch latency / eager-mode kernel gaps. We approximate as616        `(1 - gpu_util) * (1 - host_busy)` × step_time.617 618      comm_excess619        Sum of all collective-kernel duration (anything matching the rccl620        pattern). Treated as 100% recoverable elsewhere — the rule decides621        recovery_fraction.622 623      memory_headroom624        `(192 - hbm_peak) / 192 × small_constant`. Only counts the headroom625        that exceeded a healthy 70% target — running at 60% of HBM costs us626        roughly 0.07 of step time worth of optimisation surface.627 628      useful_gpu629        Whatever step time is left. Will be roughly the busy GPU time minus630        comm_excess and the kernel-shape penalty.631    """632    step_t = torch_summary.step_time_seconds or 0.0633    gpu_util = (smi.gpu_util_pct or 0.0) / 100.0  # 0..1634    host_busy = torch_summary.host_busy_fraction or 0.0635    if step_t <= 0:636        return WasteBudget()637 638    # data_wait: GPU under-utilised AND host busy → dataloader bottleneck.639    if gpu_util < 0.30 and host_busy > 0.5:640        data_wait = step_t * (1.0 - gpu_util) * host_busy641    else:642        data_wait = 0.0643 644    # host_gap: GPU idle while host idle too → launch latency / kernel gaps.645    host_gap = step_t * (1.0 - gpu_util) * (1.0 - host_busy)646 647    # comm_excess: total time in collective kernels (in seconds).648    comm_excess_ns = sum(k.duration_ns for k in kernels if k.is_collective)649    comm_excess = comm_excess_ns / 1e9650 651    # kernel_shape: fraction of GEMM time spent on un-tuned/generic kernels.652    gemm_total_ns = sum(k.duration_ns for k in kernels if k.is_gemm) or 1653    generic_gemm_ns = sum(k.duration_ns for k in kernels if k.is_generic_gemm)654    kernel_shape = (generic_gemm_ns / gemm_total_ns) * step_t * gpu_util  # cap at "real" GPU time655 656    # precision_path: only meaningful if config is fp16/fp32. Estimate from657    # kernel name tags. If config is bf16+ we leave as 0 — already optimal.658    precision_path = 0.0659    if config is None or config.precision in {"fp16", "fp32"}:660        fp16_ns = sum(k.duration_ns for k in kernels if _FP16_PATTERN.search(k.name))661        bf16_ns = sum(k.duration_ns for k in kernels if _BF16_PATTERN.search(k.name))662        denom = fp16_ns + bf16_ns663        if denom > 0:664            # Fraction of compute still on fp16. On MI300X bf16 is faster +665            # more numerically stable, so this fraction is the recoverable666            # precision_path surface.667            precision_path = (fp16_ns / denom) * step_t * gpu_util * 0.10668 669    # memory_headroom: headroom past the 70% healthy target × small constant.670    memory_headroom = 0.0671    hbm_peak = smi.hbm_peak_gb672    if hbm_peak is not None and hbm_peak > 0:673        utilisation = hbm_peak / _HBM_TOTAL_GB674        if utilisation < _HBM_HEALTHY_TARGET:675            slack = (_HBM_HEALTHY_TARGET - utilisation) / _HBM_HEALTHY_TARGET676            # Small constant: HBM slack is real but only enables a fraction of677            # potential gain (you still need a larger batch to use it). 0.05678            # of step time per "unit of slack" is a conservative anchor.679            memory_headroom = slack * step_t * 0.05680 681    # useful_gpu: everything else. Clamp to >= 0.682    spent = data_wait + host_gap + comm_excess + kernel_shape + precision_path + memory_headroom683    useful_gpu = max(0.0, step_t - spent)684 685    return WasteBudget(686        useful_gpu=useful_gpu,687        data_wait=data_wait,688        host_gap=host_gap,689        comm_excess=comm_excess,690        memory_headroom=memory_headroom,691        precision_path=precision_path,692        kernel_shape=kernel_shape,693    )694