diff --git a/src/hyperloom/agents/kernel/tests/test_bypass_classify.py b/src/hyperloom/agents/kernel/tests/test_bypass_classify.py
new file mode 100644
index 000000000..b7e802e86
--- /dev/null
+++ b/src/hyperloom/agents/kernel/tests/test_bypass_classify.py
@@ -0,0 +1,114 @@
+###############################################################################
+# SPDX-FileCopyrightText: 2026 Advanced Micro Devices, Inc.
+# SPDX-License-Identifier: MIT
+#
+# See LICENSE for license information.
+###############################################################################
+
+"""Unit tests for :mod:`_bypass_classify.classify_kernel` taxonomy coverage."""
+
+from __future__ import annotations
+
+import sys
+from pathlib import Path
+
+import pytest
+
+sys.path.insert(0, str(Path(__file__).resolve().parents[1] / "tools"))
+
+from _bypass_classify import classify_kernel  # noqa: E402
+
+
+# (device_kernel_name, expected_category) covering real sglang+aiter kernels.
+_CATEGORY_CASES = [
+    # Communication / collective.
+    ("cross_device_reduce_2stage", "Communication"),
+    ("sglang::outplace_all_reduce", "Communication"),
+    ("all_reduce", "Communication"),
+    ("allgather_vec", "Communication"),
+    ("reg_all_gather", "Communication"),
+    ("_ZN5aiter26cross_device_reduce_2stageIDF16bLi8ELb0EEEvPNS_8RankDataE", "Communication"),
+    # Attention.
+    ("_fwd_grouped_kernel_stage1", "SDPA"),
+    ("_fwd_kernel_stage2", "SDPA"),
+    ("_decode_grouped_att", "SDPA"),
+    ("mla_decode", "SDPA"),
+    ("paged_attention", "SDPA"),
+    ("_score_kernel", "SDPA"),
+    ("_combine_kernel", "SDPA"),
+    ("kda_packed_decode", "SDPA"),
+    # MoE.
+    ("fused_moe", "MoE"),
+    ("moe_gemm1", "MoE"),
+    ("moe_gemm2", "MoE"),
+    ("ck_moe", "MoE"),
+    ("fmoe", "MoE"),
+    ("moe_sorting", "MoE"),
+    ("grouped_topk", "MoE"),
+    ("moe_align", "MoE"),
+    ("topk_softmax", "MoE"),
+    ("biased_grouped_topk", "MoE"),
+    # GEMM.
+    ("gemm_a16w16", "GEMM"),
+    ("gemm_a8w8", "GEMM"),
+    ("hgemm_bf16", "GEMM"),
+    ("flatmm", "GEMM"),
+    ("Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT160x64x128", "GEMM"),
+    ("opus_gemm", "GEMM"),
+    # Normalization.
+    ("add_rmsnorm", "Normalization"),
+    ("rmsnorm", "Normalization"),
+    ("layernorm", "Normalization"),
+    ("rms_norm", "Normalization"),
+    ("add_rmsnorm_quant_kernel", "Normalization"),
+    # Elementwise.
+    ("silu_and_mul", "Elementwise"),
+    ("situ_and_mul", "Elementwise"),
+    ("act_and_mul", "Elementwise"),
+    ("elementwise_kernel", "Elementwise"),
+    ("index_elementwise", "Elementwise"),
+    ("copy_", "Elementwise"),
+    ("_to_copy", "Elementwise"),
+    ("void (anonymous namespace)::situ_and_mul_kernel", "Elementwise"),
+    # Quantization.
+    ("quant", "Quantization"),
+    ("dynamic_per_token", "Quantization"),
+    ("per_tensor_quant", "Quantization"),
+    ("static_quant_fp8", "Quantization"),
+]
+
+
+@pytest.mark.parametrize("name,expected", _CATEGORY_CASES)
+def test_category_mapping(name: str, expected: str) -> None:
+    assert classify_kernel(name).category == expected, name
+
+
+def test_communication_is_not_reusable() -> None:
+    for name in [
+        "cross_device_reduce_2stage",
+        "sglang::outplace_all_reduce",
+        "all_reduce",
+        "allgather_vec",
+        "reg_all_gather",
+        "_ZN5aiter26cross_device_reduce_2stageIDF16bLi8ELb0EEEvPNS_8RankDataE",
+    ]:
+        kc = classify_kernel(name)
+        assert kc.category == "Communication", name
+        assert kc.reusable is False, name
+
+
+def test_mangled_allreduce_via_op_name_fallback() -> None:
+    # op_name attribution present: still Communication and not reusable.
+    kc = classify_kernel("_ZN5aiter26unknown_mangledEv", op_name="sglang::outplace_all_reduce")
+    assert kc.category == "Communication"
+    assert kc.reusable is False
+
+
+def test_moe_wins_over_gemm_substring() -> None:
+    assert classify_kernel("moe_gemm1").category == "MoE"
+
+
+def test_kda_packed_decode_is_reusable_sdpa() -> None:
+    kc = classify_kernel("void (anonymous namespace)::kda_packed_decode_kernel<8, false>(x)")
+    assert kc.category == "SDPA"
+    assert kc.reusable is True
diff --git a/src/hyperloom/agents/kernel/tests/test_bypass_reader.py b/src/hyperloom/agents/kernel/tests/test_bypass_reader.py
index 2ee235707..1e399dc43 100644
--- a/src/hyperloom/agents/kernel/tests/test_bypass_reader.py
+++ b/src/hyperloom/agents/kernel/tests/test_bypass_reader.py
@@ -78,6 +78,36 @@ def test_correlation_attribution(tmp_path):
     assert kernels["paged_attention_v1"]["op_name"] == ""
 
 
+_FALLBACK_TRACE_EVENTS = [
+    # Capture-time launch of add_rmsnorm: cpu_op carries Input Dims (shape source).
+    {"cat": "cpu_op", "name": "aiter::add_rmsnorm", "args": {"External id": 100, "Input Dims": [[17, 7168]], "Input type": ["c10::BFloat16"]}},
+    {"cat": "cuda_runtime", "name": "hipLaunchKernel", "args": {"correlation": 5, "External id": 100}},
+    {"cat": "kernel", "ph": "X", "name": "add_rmsnorm_kernel", "ts": 1000, "dur": 100, "args": {"correlation": 5, "grid": [17, 1, 1], "block": [256, 1, 1]}},
+    # Graph-replay of the SAME kernel: no cpu_op link, grid-less Dispatch Task.
+    {"cat": "cuda_runtime", "name": "hipGraphLaunch", "args": {"correlation": 6}},
+    {"cat": "kernel", "ph": "X", "name": "add_rmsnorm_kernel", "ts": 2000, "dur": 100, "args": {"correlation": 6}},
+    # Pure-Triton kernel: no cpu_op, but launch carries grid/block geometry.
+    {"cat": "cuda_runtime", "name": "hipModuleLaunchKernel", "args": {"correlation": 7}},
+    {"cat": "kernel", "ph": "X", "name": "_score_kernel", "ts": 3000, "dur": 100, "args": {"correlation": 7, "grid": [17, 2, 1], "block": [512, 1, 1]}},
+]
+
+
+def test_launch_geom_and_backfill_threaded_to_rows(tmp_path):
+    payload = json.dumps({"traceEvents": _FALLBACK_TRACE_EVENTS}).encode("utf-8")
+    tf = tmp_path / "fb.trace.json"
+    tf.write_bytes(payload)
+    out = reader.analyze_trace(tf, top_k=0)
+    rows = {k["name"]: k for k in out["kernels"]}
+
+    # Same-name kernel resolved a shape at capture time -> backfill index seeded.
+    assert rows["add_rmsnorm_kernel"]["backfill_shapes"] == [[17, 7168]]
+    # Pure-Triton kernel: no cpu_op shape, but launch geometry is retained.
+    score = rows["_score_kernel"]
+    assert score["op_shapes"] == []
+    assert score["launch_grid"] == [17, 2, 1]
+    assert score["launch_block"] == [512, 1, 1]
+
+
 def test_timeline_union_math(tmp_path):
     tf = _write_trace(tmp_path / "t.trace.json")
     tl = reader.analyze_trace(tf, top_k=0)["timeline"]
diff --git a/src/hyperloom/agents/kernel/tests/test_bypass_report.py b/src/hyperloom/agents/kernel/tests/test_bypass_report.py
index 7553c2a05..8568baf59 100644
--- a/src/hyperloom/agents/kernel/tests/test_bypass_report.py
+++ b/src/hyperloom/agents/kernel/tests/test_bypass_report.py
@@ -61,6 +61,92 @@ def test_build_candidates_routing():
     assert any(c["name"] == "aten::mm" for c in cands["skipped_kernels"])
 
 
+def _one(cand_kernel):
+    """Build candidates for a single kernel row and return its candidate dict."""
+    cands = report.build_candidates(_analyze([cand_kernel]), framework="sglang", target_platform="MI300X")
+    return cands["hot_kernels"][0]
+
+
+def test_shape_provenance_torch_trace_wins():
+    # A kernel with its own cpu_op Input Dims resolves to torch_trace even when
+    # backfill/launch-grid data is also present (waterfall priority).
+    c = _one(
+        {
+            "name": "_score_kernel",
+            "op_name": "aten::score",
+            "gpu_time_us": 100.0,
+            "count": 1,
+            "op_shapes": [[17, 7168]],
+            "op_dtypes": ["c10::BFloat16"],
+            "backfill_shapes": [[99, 99]],
+            "launch_grid": [17, 2, 1],
+            "launch_block": [512, 1, 1],
+        }
+    )
+    assert c["shape_provenance"] == "torch_trace"
+    assert c["shapes"] and "17,7168" in c["shapes"][0]["shape"]
+
+
+def test_shape_provenance_capture_backfill():
+    # No own cpu_op shape, but the same-name kernel resolved a shape at capture
+    # time: inherit it, tagged capture_backfill.
+    c = _one(
+        {
+            "name": "aiter::add_rmsnorm_quant_kernel",
+            "op_name": "",
+            "gpu_time_us": 100.0,
+            "count": 1,
+            "op_shapes": [],
+            "backfill_shapes": [[17, 7168]],
+            "backfill_dtypes": ["c10::BFloat16"],
+            "launch_grid": [17, 1, 1],
+        }
+    )
+    assert c["shape_provenance"] == "capture_backfill"
+    assert c["shapes"] and "17,7168" in c["shapes"][0]["shape"]
+
+
+def test_shape_provenance_launch_grid():
+    # No cpu_op shape and no backfill: fall to launch grid/block geometry.
+    c = _one(
+        {
+            "name": "_combine_kernel",
+            "op_name": "",
+            "gpu_time_us": 100.0,
+            "count": 1,
+            "launch_grid": [17, 7, 1],
+            "launch_block": [256, 1, 1],
+        }
+    )
+    assert c["shape_provenance"] == "launch_grid"
+    s = c["shapes"][0]["shape"]
+    assert "grid=(17,7,1)" in s and "block=(256,1,1)" in s
+    # Geometry must NOT feed the analytical roofline.
+    assert c["roofline_source"] == report._RL_PLACEHOLDER
+
+
+def test_shape_provenance_tile_name():
+    # No shape and no launch geometry: parse the BLOCK_SIZE_* tile from the name.
+    c = _one(
+        {
+            "name": "_gemm_a16_w16_kernel_BLOCK_SIZE_M_32_BLOCK_SIZE_N_32_BLOCK_SIZE_K_256_x",
+            "op_name": "",
+            "gpu_time_us": 100.0,
+            "count": 1,
+        }
+    )
+    assert c["shape_provenance"] == "tile_name"
+    s = c["shapes"][0]["shape"]
+    assert "M32" in s and "N32" in s and "K256" in s
+
+
+def test_shape_provenance_unresolved_when_no_signal():
+    # Nothing to go on: stays unresolved (unchanged behaviour).
+    c = _one({"name": "mystery_kernel", "op_name": "", "gpu_time_us": 100.0, "count": 1})
+    assert c["shape_provenance"] == "unresolved"
+    assert c["shapes"] == []
+
+
 def test_build_workload_roofline_totals_covers_all_kernels():
     # 20 kernels > any top-k cap: the workload totals must sum every device
     # kernel, not just the top-k candidate list.
diff --git a/src/hyperloom/agents/kernel/tests/test_bypass_source_resolver.py b/src/hyperloom/agents/kernel/tests/test_bypass_source_resolver.py
index e619bb6df..09a7830da 100644
--- a/src/hyperloom/agents/kernel/tests/test_bypass_source_resolver.py
+++ b/src/hyperloom/agents/kernel/tests/test_bypass_source_resolver.py
@@ -14,9 +14,12 @@
 
 from __future__ import annotations
 
+import shutil
 import sys
 from pathlib import Path
 
+import pytest
+
 sys.path.insert(0, str(Path(__file__).resolve().parents[1] / "tools"))
 
 import _bypass_source_resolver as resolver  # noqa: E402
@@ -145,3 +148,53 @@ def test_editable_trace_source_rejects_generated_and_empty():
 def test_missing_json_yields_unresolved(monkeypatch):
     _patch_mapping(monkeypatch, {})
     assert resolver.resolve_source("anything", framework="vllm") == ("", "unresolved")
+
+
+def test_demangle_itanium_nested():
+    n = "_ZN5aiter26cross_device_reduce_2stageIDF16bLi8ELb0EEEvPNS_8RankDataE"
+    assert resolver._demangle_kernel_name(n) == "cross_device_reduce_2stage"
+
+
+def test_demangle_plain_triton_name_passthrough():
+    assert resolver._demangle_kernel_name("_fwd_grouped_kernel_stage1") == "_fwd_grouped_kernel_stage1"
+
+
+def test_demangle_anonymous_namespace_template():
+    n = "void (anonymous namespace)::kda_packed_decode_kernel<8, false>(x)"
+    assert resolver._demangle_kernel_name(n) == "kda_packed_decode_kernel"
+
+
+@pytest.fixture
+def repo_dir():
+    """A repo-like dir avoiding /tmp and the 'test' skip marker in its path."""
+    base = Path(__file__).resolve().parents[2] / "_bypass_repo_scan_fixture" / "src"
+    base.mkdir(parents=True, exist_ok=True)
+    yield base
+    shutil.rmtree(base.parent, ignore_errors=True)
+
+
+def test_resolve_by_kernel_name_triton_and_native(monkeypatch, repo_dir):
+    py = repo_dir / "fused.py"
+    py.write_text("@triton.jit\ndef foo(x):\n    return x\n", encoding="utf-8")
+    cu = repo_dir / "kern.cu"
+    cu.write_text("__global__ void bar(float* p) {}\n", encoding="utf-8")
+    monkeypatch.setattr(
+        resolver,
+        "_build_repo_kernel_index",
+        lambda: {"foo": str(py), "bar": str(cu), "gen": "/tmp/torchinductor_x/gen.py"},
+    )
+    assert resolver.resolve_by_kernel_name("foo") == (str(py), "repo_scan")
+    assert resolver.resolve_by_kernel_name("bar") == (str(cu), "repo_scan")
+    assert resolver.resolve_by_kernel_name("gen") == ("", "unresolved")
+    assert resolver.resolve_by_kernel_name("missing") == ("", "unresolved")
+
+
+def test_build_repo_kernel_index_scans_roots(monkeypatch, repo_dir):
+    (repo_dir / "a.py").write_text("@triton.jit\ndef tri_k(x):\n    pass\n", encoding="utf-8")
+    (repo_dir / "b.cu").write_text("__global__ void nat_k(int* p) {}\n", encoding="utf-8")
+    monkeypatch.setattr(resolver, "_repo_scan_roots", lambda: (str(repo_dir),))
+    resolver._build_repo_kernel_index.cache_clear()
+    index = resolver._build_repo_kernel_index()
+    resolver._build_repo_kernel_index.cache_clear()
+    assert index["tri_k"] == str(repo_dir / "a.py")
+    assert index["nat_k"] == str(repo_dir / "b.cu")
diff --git a/src/hyperloom/agents/kernel/tests/test_bypass_trace_analysis.py b/src/hyperloom/agents/kernel/tests/test_bypass_trace_analysis.py
index c2e577052..0eb60cd56 100644
--- a/src/hyperloom/agents/kernel/tests/test_bypass_trace_analysis.py
+++ b/src/hyperloom/agents/kernel/tests/test_bypass_trace_analysis.py
@@ -791,3 +791,70 @@ def test_discover_capture_shards_dedups_tp_ranks(tmp_path):
     labels = sorted(lbl for _p, lbl, _m in shards)
     assert labels == ["bs_16", "bs_64"]  # bs_16 deduped 2 ranks -> 1
     assert len(shards) == 2
+
+
+# ── CUDA/HIP graph-mode under-recording ──────────────────────────────────────
+
+
+def _graph_under_recorded_events():
+    """Synthetic graph-mode trace: 4 graph launches but only 1 replay's kernels
+    recorded, spanning a large wall clock (busy fraction << 0.5)."""
+    events = [{"cat": "cpu_op", "name": "aten::mm", "args": {"External id": 200}}]
+    # Four graph-launch runtime events (no External id, only correlation).
+    for i, corr in enumerate((5, 6, 7, 8)):
+        events.append(
+            {"cat": "cuda_runtime", "name": "hipGraphLaunch", "args": {"correlation": corr, "External id": None}}
+        )
+    # A normally-launched kernel (attributable) and graph-internal kernels for a
+    # single recorded replay (correlation 5). Total wall span ~10s, busy ~200us.
+    events += [
+        {"cat": "cuda_runtime", "name": "hipLaunchKernel", "args": {"correlation": 99, "External id": 200}},
+        {"cat": "kernel", "ph": "X", "name": "Cijk_Alik_Bljk_HHS", "ts": 1000, "dur": 100, "args": {"correlation": 99}},
+        {"cat": "kernel", "ph": "X", "name": "graph_fused_attn", "ts": 1100, "dur": 100, "args": {"correlation": 5}},
+        {"cat": "kernel", "ph": "X", "name": "graph_fused_mlp", "ts": 10_000_000, "dur": 100, "args": {"correlation": 5}},
+    ]
+    return events
+
+
+def test_reader_marks_graph_attributed_not_unlinked(tmp_path):
+    # Graph-launch correlations classify replayed kernels as graph-attributed,
+    # never as (unlinked); flags graph_mode / graph_under_recorded.
+    trace = tmp_path / "g.trace.json"
+    trace.write_bytes(json.dumps({"traceEvents": _graph_under_recorded_events()}).encode("utf-8"))
+    out = bta._reader.analyze_trace(str(trace), top_k=0)
+    attr = out["attribution"]
+    cov = out["graph_coverage"]
+    assert attr["graph_mode"] is True
+    assert attr["graph_launch_count"] == 4
+    assert attr["graph_attributed_kernels"] == 2
+    assert attr["unlinked_kernels"] == 0
+    assert cov["graph_under_recorded"] is True
+    assert cov["busy_fraction"] < 0.5
+    # graph-internal kernels are not surfaced as a (unlinked) op bucket
+    assert "(unlinked)" not in {o["name"] for o in out["ops"]}
+
+
+def test_reader_no_graph_mode_when_no_graph_launches(tmp_path):
+    trace = tmp_path / "ng.trace.json"
+    trace.write_bytes(json.dumps({"traceEvents": _TRACE_EVENTS}).encode("utf-8"))
+    out = bta._reader.analyze_trace(str(trace), top_k=0)
+    assert out["attribution"]["graph_mode"] is False
+    assert out["graph_coverage"]["graph_under_recorded"] is False
+
+
+def test_graph_under_recorded_skips_idle_gate_keeps_candidates(tmp_path, capsys, monkeypatch):
+    # Under-recorded graph trace: idle% is ~99% but the idle gate must NOT clear
+    # candidates; instead a bypass_graph_under_recorded warning is surfaced.
+    monkeypatch.delenv("HYPERLOOM_BYPASS_STEADY_STATE", raising=False)
+    monkeypatch.delenv("HYPERLOOM_TRACELENS_IDLE_PCT_THRESHOLD", raising=False)
+    trace = tmp_path / "g.trace.json"
+    trace.write_bytes(json.dumps({"traceEvents": _graph_under_recorded_events()}).encode("utf-8"))
+    rc, result, _ = _run(_base_argv(tmp_path, str(trace)), capsys)
+    assert rc == 0
+    assert result["timeline"]["idle_pct"] > 80.0
+    assert result["graph_coverage"]["graph_under_recorded"] is True
+    codes = {w["code"] for w in result["trace_health_warnings"]}
+    assert "bypass_graph_under_recorded" in codes
+    assert "high_gpu_idle_pct" not in codes
+    # candidates are preserved (ranked by recorded-kernel GPU share)
+    assert result["hot_kernels"], "candidates must survive the idle gate under graph under-recording"
diff --git a/src/hyperloom/agents/kernel/tools/_bypass_classify.py b/src/hyperloom/agents/kernel/tools/_bypass_classify.py
index 52af92928..a559ca772 100644
--- a/src/hyperloom/agents/kernel/tools/_bypass_classify.py
+++ b/src/hyperloom/agents/kernel/tools/_bypass_classify.py
@@ -30,14 +30,38 @@
 _RULES: list[tuple[re.Pattern, str, int]] = [
     # MemCpy (highest; also detected via Kineto cat).
     (re.compile(r"(?i)memcpy|memset"), "MemCpy", 30),
+    # Communication / collective (before generic GEMM/Elementwise/reduce). Match
+    # on the substring present inside mangled device names (e.g. aiter all-reduce).
+    (
+        re.compile(
+            r"(?i)cross_device_reduce|outplace_all_reduce|all_reduce|allreduce|"
+            r"allgather|all_gather|reg_all_gather|reduce_scatter|all_to_all|alltoall"
+        ),
+        "Communication",
+        28,
+    ),
     # MoE (specific, before generic GEMM/Elementwise).
     (re.compile(r"(?i)swiglu|fmoe|kernel_moe_gemm|MoeGemmBlockScale"), "MoE", 25),
-    (re.compile(r"(?i)moe_sorting|moe_align|topk_softmax|topkGatingSoftmax|routing"), "MoE", 22),
+    (
+        re.compile(
+            r"(?i)fused_moe|moe_gemm[12]|ck_moe|moe_sorting|moe_align|"
+            r"topk_softmax|topkGatingSoftmax|grouped_topk|biased_grouped_topk|routing"
+        ),
+        "MoE",
+        22,
+    ),
     # Attention / SDPA.
-    (re.compile(r"(?i)paged_attention|PagedAttention"), "SDPA", 20),
+    (re.compile(r"(?i)paged_attention|PagedAttention|mla_decode"), "SDPA", 20),
     (re.compile(r"(?i)attention_[23]d|unified_attention"), "SDPA", 20),
     (re.compile(r"(?i)flash_attn|flash_fwd|fmha"), "SDPA", 20),
-    (re.compile(r"(?i)_fwd_kernel|reduce_segments"), "SDPA", 18),
+    (
+        re.compile(
+            r"(?i)_fwd_kernel|_fwd_grouped_kernel|_decode_grouped_att|"
+            r"_score_kernel|_combine_kernel|kda_packed_decode|reduce_segments"
+        ),
+        "SDPA",
+        18,
+    ),
     # KV cache store.
     (re.compile(r"(?i)reshape_and_cache|concat_and_cache"), "KVCacheStore", 20),
     # Normalization.
@@ -56,14 +80,31 @@
         "Quantization",
         18,
     ),
-    (re.compile(r"(?i)\bquant\b|quantize"), "Quantization", 6),
+    (
+        re.compile(r"(?i)dynamic_per_token|per_token.*quant|static_quant_fp8|\bquant\b|quantize"),
+        "Quantization",
+        6,
+    ),
     # Activation / elementwise.
-    (re.compile(r"(?i)silu|swish|\bgelu\b|act_and_mul"), "Elementwise", 15),
+    (re.compile(r"(?i)silu|swish|\bgelu\b|silu_and_mul|situ_and_mul|act_and_mul"), "Elementwise", 15),
     (re.compile(r"(?i)embedding|gather_kernel|vectorized_gather"), "Elementwise", 14),
-    (re.compile(r"(?i)elementwise|CatArrayBatchedCopy|direct_copy_kernel|_to_copy\b|FillFunctor"), "Elementwise", 10),
+    (
+        re.compile(
+            r"(?i)elementwise|index_elementwise|CatArrayBatchedCopy|direct_copy_kernel|"
+            r"copy_|_to_copy\b|FillFunctor"
+        ),
+        "Elementwise",
+        10,
+    ),
     # GEMM (vendor + generic).
     (re.compile(r"(?i)Cijk_|wvSplitK|splitKreduce|hipblaslt|rocblas|cublas|nvjet"), "GEMM", 12),
-    (re.compile(r"(?i)kernel_gemm_xdl_cshuffle|_gemm_a\d+w\d+|_gemm_a16_w16"), "GEMM", 12),
+    (
+        re.compile(
+            r"(?i)kernel_gemm_xdl_cshuffle|_?gemm_a\d+w\d+|_gemm_a16_w16|hgemm_bf16|flatmm|opus_gemm"
+        ),
+        "GEMM",
+        12,
+    ),
     (re.compile(r"(?i)scaled_mm|\bgemm\b|matmul|\bbmm\b"), "GEMM", 8),
     # Triton / generic catch-alls (lowest).
     (re.compile(r"(?i)triton_poi_fused|triton_red_fused|triton_per_fused"), "Elementwise", 3),
@@ -75,8 +116,10 @@
 _VENDOR_BINARY_RE = re.compile(r"(?i)Cijk_|wvSplitK|splitKreduce|hipblaslt|rocblas|cublas|nvjet_tst|miopen|cudnn")
 # Native-source kernels that are rewritable (triton / aiter / CK / vLLM native).
 _REUSABLE_RE = re.compile(
-    r"(?i)triton_|^_fwd_kernel|aiter|ck_tile|paged_attention|reshape_and_cache|"
-    r"rmsnorm|rms_norm|add_rmsnorm|silu|per_tensor_quant|scaled_quant|data_to_scale|initializeScale|"
+    r"(?i)triton_|^_fwd_kernel|_fwd_grouped_kernel|_decode_grouped_att|_score_kernel|_combine_kernel|"
+    r"kda_packed_decode|aiter|ck_tile|paged_attention|mla_decode|reshape_and_cache|"
+    r"rmsnorm|rms_norm|add_rmsnorm|silu|situ_and_mul|act_and_mul|"
+    r"per_tensor_quant|dynamic_per_token|static_quant_fp8|scaled_quant|data_to_scale|initializeScale|"
     r"fusedlnmodulate|ada_?ln|modulate|scale_shift"
 )
 
@@ -86,6 +129,13 @@
 # row, so rows MUST stay ordered most-specific-first. Norm keeps only explicit
 # ``*_norm`` names so a bare ``\bnorm\b`` does not mis-map ``aten::norm``.
 _OP_RULES: list[tuple[re.Pattern, str]] = [
+    (
+        re.compile(
+            r"(?i)cross_device_reduce|outplace_all_reduce|all_reduce|allreduce|"
+            r"allgather|all_gather|reg_all_gather|reduce_scatter|all_to_all|alltoall"
+        ),
+        "Communication",
+    ),
     (re.compile(r"(?i)convolution|conv[123]d|conv_transpose|_convolution"), "Convolution"),
     (
         re.compile(
@@ -178,6 +228,9 @@ def classify_kernel(name: str, *, gpu_cat: str = "", op_name: str = "") -> Kerne
 
     if category == "MemCpy":
         return KernelClass("MemCpy", False, "device memcpy/memset (not a rewritable kernel)")
+    # Collectives are runtime/library primitives, never a rewritable candidate.
+    if category == "Communication":
+        return KernelClass(category, False, "collective/communication primitive (not a rewritable kernel)")
     if _VENDOR_BINARY_RE.search(n):
         return KernelClass(category, False, "vendor backend library (precompiled binary, no rewritable source)")
     if _REUSABLE_RE.search(n):
diff --git a/src/hyperloom/agents/kernel/tools/_bypass_report.py b/src/hyperloom/agents/kernel/tools/_bypass_report.py
index 73bcbcfc7..4c34a7283 100644
--- a/src/hyperloom/agents/kernel/tools/_bypass_report.py
+++ b/src/hyperloom/agents/kernel/tools/_bypass_report.py
@@ -29,7 +29,7 @@
 from _analysis_md import render_report
 from _bypass_roofline import compute_roofline
 from _kernel_category import canonical_category
-from _bypass_source_resolver import editable_trace_source, resolve_source
+from _bypass_source_resolver import editable_trace_source, resolve_by_kernel_name, resolve_source
 from _idle_gate import resolve_idle_pct_threshold
 from _roofline_source import PLACEHOLDER as _RL_PLACEHOLDER
 from _task_group_contract import (
@@ -147,6 +147,54 @@ def _trace_shape_entries(op_shapes: Any, op_dtypes: Any, call_count: int) -> lis
     return [{"call_num": int(call_count) if call_count else 1, "shape": "<br>".join(operands)}]
 
 
+# Triton autotune names embed the tile shape, e.g.
+# ``_gemm_a16_w16_kernel_BLOCK_SIZE_M_32_BLOCK_SIZE_N_32_BLOCK_SIZE_K_256_...``.
+_TILE_NAME_RE = re.compile(r"BLOCK_SIZE_([A-Za-z]+)_(\d+)")
+
+
+def _launch_grid_shape_entries(grid: Any, block: Any, call_count: int) -> list[dict[str, Any]]:
+    """Build a shape entry from a kernel's launch geometry (grid/block).
+
+    Fallback for kernels whose correlation->cpu_op chain is broken (Triton
+    direct-launch, graph replay): the launch grid/block is not the operand
+    tensor shape but a coarse, dispatchable geometry. Returns ``[]`` when no
+    positive dimension is present.
+    """
+    def _dims(v: Any) -> list[int]:
+        out: list[int] = []
+        for d in v if isinstance(v, (list, tuple)) else []:
+            try:
+                n = int(d)
+            except (TypeError, ValueError):
+                return []
+            out.append(n)
+        return out
+
+    g = _dims(grid)
+    b = _dims(block)
+    parts: list[str] = []
+    if any(x > 0 for x in g):
+        parts.append("grid=(" + ",".join(str(x) for x in g) + ")")
+    if any(x > 0 for x in b):
+        parts.append("block=(" + ",".join(str(x) for x in b) + ")")
+    if not parts:
+        return []
+    return [{"call_num": int(call_count) if call_count else 1, "shape": "<br>".join(parts)}]
+
+
+def _tile_name_shape_entries(kernel_name: str, call_count: int) -> list[dict[str, Any]]:
+    """Extract a tile shape from a Triton autotune kernel name (``BLOCK_SIZE_*``).
+
+    Lowest-priority fallback: yields a ``M32<br>N32<br>K256``-style tile shape
+    when the name encodes it. Returns ``[]`` when no tile token is present.
+    """
+    matches = _TILE_NAME_RE.findall(kernel_name or "")
+    if not matches:
+        return []
+    body = "<br>".join(f"{axis}{val}" for axis, val in matches)
+    return [{"call_num": int(call_count) if call_count else 1, "shape": body}]
+
+
 def _source_type_for_op(op_name: str) -> str:
     """Best-effort source-type guess from a launching op name.
 
@@ -361,19 +409,47 @@ def build_candidates(
         display = op_name or _short_name(kname)
 
         # Source resolution. Priority: (1) a Triton kernel_file from the trace's
-        # cpu_op args; (2) op_to_source.json lookup; else unresolved.
+        # cpu_op args; (2) op_to_source.json lookup; (3) repo-scan by device
+        # kernel name; else unresolved.
         source_file = editable_trace_source(k.get("op_kernel_file", "") or "", k.get("op_kernel_backend", "") or "")
         source_method = "trace_kernel_file" if source_file else "unresolved"
         if not source_file and op_name:
             source_file, method = resolve_source(op_name, framework=framework, device_kernel_name=kname)
             if source_file:
                 source_method = method
+        if not source_file and kname:
+            source_file, method = resolve_by_kernel_name(kname)
+            if source_file:
+                source_method = method
 
-        # Real per-arg dims/dtypes from the trace, rendered into the downstream
-        # shape-string contract.
+        # Shape resolution waterfall (provenance records the source):
+        #   1. torch_trace      -- this kernel's own cpu_op Input Dims (precise)
+        #   2. capture_backfill -- same-name kernel's capture-time shape
+        #   3. launch_grid      -- this kernel's launch grid/block geometry
+        #   4. tile_name        -- BLOCK_SIZE_* tile embedded in the kernel name
+        #   5. unresolved       -- none of the above
+        _count = k.get("count") or 0
         op_shapes = k.get("op_shapes") or []
         op_dtypes = k.get("op_dtypes") or []
-        shape_entries = _trace_shape_entries(op_shapes, op_dtypes, k.get("count") or 0)
+        shape_entries = _trace_shape_entries(op_shapes, op_dtypes, _count)
+        shape_provenance = "torch_trace" if shape_entries else ""
+        if not shape_entries:
+            bf_shapes = k.get("backfill_shapes") or []
+            bf_dtypes = k.get("backfill_dtypes") or []
+            shape_entries = _trace_shape_entries(bf_shapes, bf_dtypes, _count)
+            if shape_entries:
+                op_dtypes = bf_dtypes
+                shape_provenance = "capture_backfill"
+        if not shape_entries:
+            shape_entries = _launch_grid_shape_entries(k.get("launch_grid"), k.get("launch_block"), _count)
+            if shape_entries:
+                shape_provenance = "launch_grid"
+        if not shape_entries:
+            shape_entries = _tile_name_shape_entries(kname, _count)
+            if shape_entries:
+                shape_provenance = "tile_name"
+        if not shape_provenance:
+            shape_provenance = "unresolved"
 
         # Benchmark discovery is opt-in; a routable kernel's on-disk
         # test/benchmark can seed downstream harness generation.
@@ -427,14 +503,22 @@ def build_candidates(
             "shapes": shape_entries,
             "input_shapes": shape_entries,
             "input_dtypes": op_dtypes,
-            "shape_provenance": "torch_trace" if shape_entries else "unresolved",
+            "shape_provenance": shape_provenance,
         }
         # Analytical roofline: derive bound_type / AI / efficiency from captured
         # shapes + measured time for EVERY estimable kernel (rocprof enrichment
-        # later refines it to a measured roofline).
+        # later refines it to a measured roofline). Only precise operand shapes
+        # (torch_trace / capture_backfill) feed the AI math; launch-grid and
+        # tile-name fallbacks are geometry, not operand dims, so they would
+        # poison the roofline -- skip analytical estimation for them.
+        _roofline_shape = (
+            shape_entries[0]["shape"]
+            if shape_entries and shape_provenance in ("torch_trace", "capture_backfill")
+            else ""
+        )
         rl = compute_roofline(
             category=kc.category,
-            shape_str=shape_entries[0]["shape"] if shape_entries else "",
+            shape_str=_roofline_shape,
             gpu_time_us=float(cand["duration_us"] or 0.0),
             call_count=int(cand["call_count"] or 1),
             gpu_type=target_platform,
diff --git a/src/hyperloom/agents/kernel/tools/_bypass_source_resolver.py b/src/hyperloom/agents/kernel/tools/_bypass_source_resolver.py
index bf8b97b97..f6ed4e1bb 100644
--- a/src/hyperloom/agents/kernel/tools/_bypass_source_resolver.py
+++ b/src/hyperloom/agents/kernel/tools/_bypass_source_resolver.py
@@ -237,6 +237,124 @@ def resolve_source(
     return sources[0], "op_to_source"
 
 
+# Length-prefixed token of an Itanium name (e.g. "5aiter" -> "aiter").
+_ITANIUM_TOKEN_RE = re.compile(r"(\d+)")
+# Triton kernel definition: @triton.jit then optional decorators then def NAME.
+_TRITON_DEF_RE = re.compile(r"@triton\.jit[^\n]*\n(?:\s*@[^\n]*\n)*\s*def\s+(\w+)")
+# Native kernel definition: __global__ with optional qualifiers then NAME.
+_GLOBAL_DEF_RE = re.compile(
+    r"__global__\s*(?:(?:void|static|inline|__forceinline__|"
+    r"__launch_bounds__\s*\([^)]*\))\s*)*(\w+)\s*[\(<]"
+)
+# Directories/paths to skip while scanning source repos.
+_SCAN_SKIP_MARKERS = ("/__pycache__", "/3rdparty/", "/example", "/test", "/jit/build/", "/.git/")
+_TRITON_SCAN_EXTS = (".py",)
+_NATIVE_SCAN_EXTS = (".cu", ".cuh", ".hip", ".h")
+
+
+def _demangle_kernel_name(name: str) -> str | None:
+    """Extract the bare function identifier from a device kernel name.
+
+    Handles Itanium mangling (``_ZN<len><ns>...`` / ``_Z<len><name>...``) and
+    plain C++/Triton names (strips ``void``, namespaces, and template/arg tails).
+    """
+    n = (name or "").strip()
+    if not n:
+        return None
+    if n.startswith("_ZN") or n.startswith("_Z"):
+        body = n[3:] if n.startswith("_ZN") else n[2:]
+        tokens: list[str] = []
+        i = 0
+        while i < len(body) and body[i].isdigit():
+            j = i
+            while j < len(body) and body[j].isdigit():
+                j += 1
+            length = int(body[i:j])
+            token = body[j : j + length]
+            if not token:
+                break
+            tokens.append(token)
+            i = j + length
+        if tokens:
+            return tokens[-1] if n.startswith("_ZN") else tokens[0]
+        return None
+    if n.startswith("void "):
+        n = n[len("void ") :].strip()
+    n = n.replace("(anonymous namespace)::", "")
+    n = re.sub(r"<.*$", "", n)
+    n = re.sub(r"\(.*$", "", n)
+    n = n.strip()
+    if "::" in n:
+        n = n.rsplit("::", 1)[-1]
+    return n or None
+
+
+@functools.lru_cache(maxsize=1)
+def _repo_scan_roots() -> tuple[str, ...]:
+    """Discover live sglang/aiter source roots (and aiter csrc) without importing."""
+    roots: list[str] = []
+    seen: set[str] = set()
+    for pkg in ("sglang", "aiter"):
+        try:
+            spec = importlib.util.find_spec(pkg)
+        except (ImportError, ValueError, ModuleNotFoundError):
+            continue
+        if spec is None:
+            continue
+        for loc in list(getattr(spec, "submodule_search_locations", None) or []):
+            for cand in (loc, os.path.join(os.path.dirname(loc), "csrc")):
+                if cand not in seen and os.path.isdir(cand):
+                    seen.add(cand)
+                    roots.append(cand)
+    return tuple(roots)
+
+
+@functools.lru_cache(maxsize=1)
+def _build_repo_kernel_index() -> dict[str, str]:
+    """Map kernel function name -> source path by scanning repo roots once."""
+    index: dict[str, str] = {}
+    for root in _repo_scan_roots():
+        for dirpath, _dirs, files in os.walk(root):
+            if any(m in dirpath for m in _SCAN_SKIP_MARKERS):
+                continue
+            for fname in files:
+                path = os.path.join(dirpath, fname)
+                if any(m in path for m in _SCAN_SKIP_MARKERS):
+                    continue
+                low = fname.lower()
+                if low.endswith(_TRITON_SCAN_EXTS):
+                    pattern, marker = _TRITON_DEF_RE, "@triton.jit"
+                elif low.endswith(_NATIVE_SCAN_EXTS):
+                    pattern, marker = _GLOBAL_DEF_RE, "__global__"
+                else:
+                    continue
+                try:
+                    with open(path, encoding="utf-8", errors="ignore") as fh:
+                        text = fh.read()
+                except OSError:
+                    continue
+                if marker not in text:
+                    continue
+                for match in pattern.finditer(text):
+                    index.setdefault(match.group(1), path)
+    return index
+
+
+def resolve_by_kernel_name(device_kernel_name: str) -> tuple[str, str]:
+    """Resolve a device kernel name to an editable source via repo scan.
+
+    Demangles the kernel name and looks it up in the repo kernel index, returning
+    ``(path, "repo_scan")`` on an editable on-disk hit, else ``("", "unresolved")``.
+    """
+    bare = _demangle_kernel_name(device_kernel_name)
+    if not bare:
+        return "", "unresolved"
+    path = _build_repo_kernel_index().get(bare)
+    if path and is_editable_source(path) and _exists(path):
+        return path, "repo_scan"
+    return "", "unresolved"
+
+
 def editable_trace_source(kernel_file: str, kernel_kind: str = "") -> str:
     """Return a trace-provided Triton ``kernel_file`` iff it is an editable source.
 
@@ -257,4 +375,4 @@ def editable_trace_source(kernel_file: str, kernel_kind: str = "") -> str:
     return kf if is_editable_source(kf, kernel_kind or None) else ""
 
 
-__all__ = ["resolve_source", "editable_trace_source", "is_editable_source"]
+__all__ = ["resolve_source", "resolve_by_kernel_name", "editable_trace_source", "is_editable_source"]
diff --git a/src/hyperloom/agents/kernel/tools/_bypass_trace_reader.py b/src/hyperloom/agents/kernel/tools/_bypass_trace_reader.py
index 3d5ad2a3a..17cb745d7 100644
--- a/src/hyperloom/agents/kernel/tools/_bypass_trace_reader.py
+++ b/src/hyperloom/agents/kernel/tools/_bypass_trace_reader.py
@@ -344,6 +344,9 @@ def _finalize(
     window: tuple[float, float] | None,
     top_k: int,
     emit_launches: bool = False,
+    graph_launch_corrs: frozenset[int] | None = None,
+    graph_launch_count: int = 0,
+    corr_to_launch_geom: dict[Any, tuple[Any, Any]] | None = None,
 ) -> dict[str, Any]:
     """Build timeline + op/kernel aggregates from buffered device events.
 
@@ -418,12 +421,30 @@ def _clip(a: float, b: float) -> tuple[float, float] | None:
     attributed_kernels = 0
     unlinked_us = 0.0
     unlinked_kernels = 0
+    # Graph-internal kernels launched via a captured CUDA/HIP graph: they carry a
+    # correlation matching a graph-launch runtime event (which has no External id),
+    # so they resolve to no cpu_op but are not a genuine attribution failure.
+    graph_corrs = graph_launch_corrs or frozenset()
+    graph_kernels = 0
+    graph_gpu_us = 0.0
+    geom = corr_to_launch_geom or {}
+    # Name-keyed shape backfill: a kernel replayed inside a CUDA graph loses its
+    # correlation->cpu_op link, but the SAME kernel resolved a shape on an
+    # eager/capture-time launch. Record the first shape-carrying meta per kernel
+    # name so a later shapeless launch can inherit it (provenance downgrades).
+    kern_name_backfill_meta: dict[str, dict[str, Any]] = {}
+    # Name-keyed launch geometry (grid/block), first-seen per kernel name.
+    kern_name_launch_geom: dict[str, tuple[Any, Any]] = {}
     for name, dur, corr, _ts, _end in k_events:
         extid = corr_to_extid.get(corr) if corr is not None else None
         op_name = extid_to_opname.get(extid) if extid is not None else None
         if op_name:
             attributed_us += dur
             attributed_kernels += 1
+        elif corr is not None and corr in graph_corrs:
+            op_name = "(graph)"
+            graph_kernels += 1
+            graph_gpu_us += dur
         else:
             op_name = "(unlinked)"
             unlinked_us += dur
@@ -440,6 +461,13 @@ def _clip(a: float, b: float) -> tuple[float, float] | None:
             meta = extid_to_opmeta.get(extid)
             if meta:
                 kern_op_meta.setdefault(name, {}).setdefault(op_name, meta)
+                # Seed the name-keyed backfill from any shape-carrying meta.
+                if meta.get("shapes") and name not in kern_name_backfill_meta:
+                    kern_name_backfill_meta[name] = meta
+        if name not in kern_name_launch_geom:
+            g = geom.get(corr)
+            if g is not None and (g[0] is not None or g[1] is not None):
+                kern_name_launch_geom[name] = g
 
     def _majority_op(kernel_name: str) -> str:
         """Return the highest-GPU-time real launching op for a kernel name.
@@ -448,7 +476,7 @@ def _majority_op(kernel_name: str) -> str:
         """
         best, best_dur = "", 0.0
         for op, d in (kern_op.get(kernel_name) or {}).items():
-            if op != "(unlinked)" and d > best_dur:
+            if op not in ("(unlinked)", "(graph)") and d > best_dur:
                 best, best_dur = op, d
         return best
 
@@ -471,6 +499,16 @@ def _majority_op_meta(kernel_name: str) -> dict[str, Any]:
         total_ms = 0.0
     idle_ms = max(0.0, total_ms - busy_ms)
 
+    # Graph coverage health: under continuous graph replay roctracer's activity
+    # buffer overflows, so only ~1 replay's kernels are recorded. Flag when
+    # graphs are replaying yet recorded busy covers <50% of the wall span, or
+    # many launches map to a single recorded replay's worth of kernels.
+    graph_mode = graph_launch_count > 0
+    busy_fraction = round(busy_ms / total_ms, 4) if total_ms > 0 else 0.0
+    graph_under_recorded = graph_mode and (
+        busy_fraction < 0.5 or (graph_launch_count >= 4 and graph_kernels > 0 and busy_fraction < 0.9)
+    )
+
     gpu_kernel_total_us = sum(a.dur_us for a in kern_agg.values())
 
     def _top(agg: dict[str, _Agg], denom_us: float, *, attach_op: bool = False) -> list[dict[str, Any]]:
@@ -490,6 +528,14 @@ def _top(agg: dict[str, _Agg], denom_us: float, *, attach_op: bool = False) -> l
                 row["op_dtypes"] = meta.get("dtypes") or []
                 row["op_kernel_file"] = meta.get("kernel_file") or ""
                 row["op_kernel_backend"] = meta.get("kernel_backend") or ""
+                # Shape fallbacks for a kernel whose own launch had no cpu_op
+                # shape: (a) same-name capture-time shape, (b) launch geometry.
+                bf = kern_name_backfill_meta.get(nm) or {}
+                row["backfill_shapes"] = bf.get("shapes") or []
+                row["backfill_dtypes"] = bf.get("dtypes") or []
+                lg = kern_name_launch_geom.get(nm)
+                row["launch_grid"] = list(lg[0]) if lg and lg[0] is not None else []
+                row["launch_block"] = list(lg[1]) if lg and lg[1] is not None else []
             rows.append(row)
         rows.sort(key=lambda r: r["gpu_time_ms"], reverse=True)
         return rows if top_k is None or top_k <= 0 else rows[:top_k]
@@ -547,6 +593,17 @@ def _top(agg: dict[str, _Agg], denom_us: float, *, attach_op: bool = False) -> l
             "cuda_runtime_links": len(corr_to_extid),
             "cpu_ops": len(extid_to_opname),
             "gpu_memcpy_count": memcpy_count,
+            "graph_mode": graph_mode,
+            "graph_launch_count": graph_launch_count,
+            "graph_attributed_kernels": graph_kernels,
+            "graph_attributed_gpu_ms": round(graph_gpu_us / 1000.0, 4),
+        },
+        "graph_coverage": {
+            "graph_mode": graph_mode,
+            "graph_launch_count": graph_launch_count,
+            "graph_attributed_kernels": graph_kernels,
+            "busy_fraction": busy_fraction,
+            "graph_under_recorded": graph_under_recorded,
         },
     }
 
@@ -577,7 +634,8 @@ def analyze_trace(
           ``timeline`` (total/busy/idle/kernel/memcpy ms),
           ``ops`` (op-level GPU-time aggregates, desc by gpu time),
           ``kernels`` (device-kernel aggregates, desc by gpu time),
-          ``attribution`` (coverage stats),
+          ``attribution`` (coverage stats + graph-mode signals),
+          ``graph_coverage`` (graph-mode / under-recording health signals),
           ``annotation_windows`` (gpu_user_annotation name/ts/dur),
           ``aggregation_scope`` (``steady_state`` or ``full_trace``),
           ``steady_window`` (window meta when steady state was applied),
@@ -599,11 +657,18 @@ def analyze_trace(
     extid_to_opname: dict[int, str] = {}
     # Compact per-op meta (first-seen) for shape + Triton-source enrichment.
     extid_to_opmeta: dict[int, dict[str, Any]] = {}
+    # Launch geometry (grid/block) per correlation, kept so a kernel whose
+    # correlation->cpu_op shape chain is broken (Triton direct-launch,
+    # graph replay) can still surface a launch-grid shape fallback.
+    corr_to_launch_geom: dict[Any, tuple[Any, Any]] = {}
     # Buffered device events so one pass serves both full-trace and steady-window
     # aggregation without re-reading the trace.
     k_events: list[tuple[str, float, Any, float, float]] = []
     m_events: list[tuple[float, float, float]] = []
     annotation_windows: list[dict[str, Any]] = []
+    # Correlations of CUDA/HIP graph-launch runtime events (no External id), so
+    # their replayed kernels are classified graph-attributed, not (unlinked).
+    graph_launch_corrs: set[int] = set()
     event_total = 0
 
     fobj = _open_trace_binary(tf)
@@ -617,6 +682,8 @@ def analyze_trace(
                 extid = args.get("External id")
                 if cid is not None and extid is not None:
                     corr_to_extid[cid] = extid
+                if cid is not None and "GraphLaunch" in (ev.get("name", "") or ""):
+                    graph_launch_corrs.add(cid)
                 continue
             if cat == "cpu_op":
                 args = ev.get("args") or {}
@@ -649,7 +716,15 @@ def analyze_trace(
                 end = ts + dur
                 name = ev.get("name", "") or ""
                 if cat == _GPU_KERNEL_CAT:
-                    corr = (ev.get("args") or {}).get("correlation")
+                    kargs = ev.get("args") or {}
+                    corr = kargs.get("correlation")
+                    # Retain launch grid/block: for Triton (hipModuleLaunchKernel)
+                    # and graph-replay kernels the correlation->cpu_op shape chain
+                    # is broken, but the launch geometry is a shape fallback.
+                    grid = kargs.get("grid")
+                    block = kargs.get("block")
+                    if grid is not None or block is not None:
+                        corr_to_launch_geom[corr] = (grid, block)
                     k_events.append((name, dur, corr, ts, end))
                 else:
                     m_events.append((dur, ts, end))
@@ -674,6 +749,9 @@ def analyze_trace(
         window=window,
         top_k=top_k,
         emit_launches=emit_launches,
+        graph_launch_corrs=frozenset(graph_launch_corrs),
+        graph_launch_count=len(graph_launch_corrs),
+        corr_to_launch_geom=corr_to_launch_geom,
     )
     body["attribution"]["annotation_window_count"] = len(annotation_windows)
 
diff --git a/src/hyperloom/agents/kernel/tools/_idle_gate.py b/src/hyperloom/agents/kernel/tools/_idle_gate.py
index 50ab53c05..0f056b226 100644
--- a/src/hyperloom/agents/kernel/tools/_idle_gate.py
+++ b/src/hyperloom/agents/kernel/tools/_idle_gate.py
@@ -82,3 +82,36 @@ def build_high_idle_warning(
             "can route to params/backends."
         ),
     }
+
+
+def build_graph_under_recorded_warning(
+    *,
+    graph_launch_count: int,
+    idle_pct: float | None = None,
+) -> dict[str, Any]:
+    """Build the ``trace_health_warnings[]`` entry for a graph under-recorded trace.
+
+    Under continuous CUDA/HIP graph replay the profiler activity buffer overflows
+    and captures only ~1 of ``graph_launch_count`` replays, so idle% is unreliable
+    and must not gate candidates; ranking by recorded-kernel GPU share stays valid.
+
+    Args:
+        graph_launch_count: Number of graph-launch runtime events in the trace.
+        idle_pct: The (unreliable) measured GPU idle percentage, for context.
+
+    Returns:
+        The structured ``bypass_graph_under_recorded`` warning entry.
+    """
+    idle_note = f" (computed idle% {idle_pct:.2f}% is unreliable here)" if isinstance(idle_pct, (int, float)) else ""
+    return {
+        "code": "bypass_graph_under_recorded",
+        "severity": "warning",
+        "graph_launch_count": graph_launch_count,
+        "message": (
+            f"graph-mode trace under-recorded: only ~1 of {graph_launch_count} graph "
+            f"replays captured (profiler activity-buffer overflow under continuous GPU "
+            f"saturation){idle_note}; idle% is unreliable and the idle gate is skipped. "
+            "Hot-kernel candidates are still ranked by recorded-kernel GPU share, which "
+            "is a representative sample of one replay."
+        ),
+    }
diff --git a/src/hyperloom/agents/kernel/tools/bypass_trace_analysis.py b/src/hyperloom/agents/kernel/tools/bypass_trace_analysis.py
index 17a6dcc88..c137b6c29 100644
--- a/src/hyperloom/agents/kernel/tools/bypass_trace_analysis.py
+++ b/src/hyperloom/agents/kernel/tools/bypass_trace_analysis.py
@@ -50,6 +50,7 @@
 except Exception:  # noqa: BLE001 — standalone invocation without the package installed.
     _shared_build_provenance = None
 from _idle_gate import (  # noqa: E402
+    build_graph_under_recorded_warning,
     build_high_idle_warning,
     resolve_idle_pct_threshold,
 )
@@ -654,7 +655,18 @@ def main(argv: list[str] | None = None) -> int:
     # surface a high_gpu_idle_pct warning for the Coordinator.
     idle_pct_value = (analyze.get("timeline") or {}).get("idle_pct")
     idle_pct_threshold = resolve_idle_pct_threshold()
-    if isinstance(idle_pct_value, (int, float)) and float(idle_pct_value) > idle_pct_threshold:
+    # Graph under-recording makes idle% unreliable: skip the idle gate (keep
+    # candidates ranked by recorded-kernel GPU share) and surface a health warning.
+    graph_coverage = analyze.get("graph_coverage") or {}
+    graph_under_recorded = bool(graph_coverage.get("graph_under_recorded"))
+    if graph_under_recorded:
+        trace_health_warnings.append(
+            build_graph_under_recorded_warning(
+                graph_launch_count=int(graph_coverage.get("graph_launch_count", 0) or 0),
+                idle_pct=float(idle_pct_value) if isinstance(idle_pct_value, (int, float)) else None,
+            )
+        )
+    elif isinstance(idle_pct_value, (int, float)) and float(idle_pct_value) > idle_pct_threshold:
         for _cand_key in ("hot_kernels", "routable_kernels", "skipped_kernels", "task_groups"):
             candidates[_cand_key] = []
         trace_health_warnings.append(
@@ -812,6 +824,7 @@ def main(argv: list[str] | None = None) -> int:
         "orchestrator_mode": "bypass",
         "timeline": analyze.get("timeline") or {},
         "attribution": analyze.get("attribution") or {},
+        "graph_coverage": analyze.get("graph_coverage") or {},
         "trace_health_warnings": trace_health_warnings,
         "artifact_paths": {
             "trace_report_path": str(analysis_md_path),
diff --git a/src/hyperloom/orchestrator/actions/executors/roofline.py b/src/hyperloom/orchestrator/actions/executors/roofline.py
index 0fc58a17c..3f475d5c5 100644
--- a/src/hyperloom/orchestrator/actions/executors/roofline.py
+++ b/src/hyperloom/orchestrator/actions/executors/roofline.py
@@ -288,8 +288,12 @@ async def __call__(self, ctx: RunnerContext) -> dict[str, Any]:
         # (profile_no_trace_failed) instead of collapsing into profile_failed.
         last_phase = "profile"
         # After a cuda-graph capture crash the next attempt boots eager so the
-        # torch-profiler stream capture cannot collide.
-        disable_cuda_graph = False
+        # torch-profiler stream capture cannot collide. Operators can arm it
+        # upfront too: the profiler cannot decompose cuda-graph replay, so a
+        # graph-mode trace reports near-zero GPU time and trips the idle gate.
+        disable_cuda_graph = os.environ.get(
+            "HYPERLOOM_PROFILE_DISABLE_CUDA_GRAPH", ""
+        ).strip().lower() in {"1", "true", "yes", "on"}
         # Resolve framework so the eager fallback picks the correct flag (vLLM
         # --enforce-eager, sglang --disable-cuda-graph).
         framework = self._resolve_framework(ctx)