diff --git a/.github/federation.json b/.github/federation.json
index 21dbeb80..8b09dc01 100644
--- a/.github/federation.json
+++ b/.github/federation.json
@@ -95,10 +95,52 @@
       "repo": "ROCm/rocm-systems",
       "branch": "develop",
       "license": "MIT",
+      "fixtures": [
+        {
+          "path": "projects/rocprofiler-compute/sample",
+          "as": "rocprof-compute-sample",
+          "workspace_prefix": "../../sample",
+          "files": [
+            "common.h",
+            "lds.hip",
+            "mat_mul_max.hip",
+            "occupancy.hip",
+            "simple_net.py",
+            "vcopy.cpp",
+            "vmem.hip",
+            "vrandom_access.cpp",
+            "vsequential_access.cpp",
+            "membw_analysis_test_suite/gl2_backpressure.hip",
+            "wmma_gemm_fp16/Makefile",
+            "wmma_gemm_fp16/main.cpp",
+            "wmma_gemm_fp16/wmma_gemm_fp16.hip",
+            "wmma_gemm_fp16/wmma_gemm_fp16.hpp",
+            "wmma_gemm_fp16/wmma_gemm_fp16_config.hpp"
+          ]
+        }
+      ],
       "skills": [
         {
           "path": "projects/rocprofiler-sdk/skills/pc-sampling",
           "as": "rocprofiler-pc-sampling"
+        },
+        {
+          "path": "projects/rocprofiler-compute/skills/rocprof-compute-kernel-bottleneck"
+        },
+        {
+          "path": "projects/rocprofiler-compute/skills/rocprof-compute-memory"
+        },
+        {
+          "path": "projects/rocprofiler-compute/skills/rocprof-compute-pc-sampling"
+        },
+        {
+          "path": "projects/rocprofiler-compute/skills/rocprof-compute-roofline"
+        },
+        {
+          "path": "projects/rocprofiler-compute/skills/rocprof-compute-speed-of-light"
+        },
+        {
+          "path": "projects/rocprofiler-compute/skills/rocprof-compute-torch-trace"
         }
       ]
     }
diff --git a/.github/scripts/federate_skills.py b/.github/scripts/federate_skills.py
index 347d3225..d7e664c4 100644
--- a/.github/scripts/federate_skills.py
+++ b/.github/scripts/federate_skills.py
@@ -32,6 +32,12 @@
    not vendored: they are left out of the copy and of the content hash, and
    named in the run log and the pull request body. The one file kept when
    upstream lacks it is `evals/machine.yml` (`LOCAL_FALLBACK_FILES`).
+   A source may also declare `fixtures`: files from a directory that is not
+   a skill (the rocprofiler-compute sample tree) copied once under
+   `skills/<as>/`. Only the listed `files` are copied.
+   Vendored `evals.json` workspace values that start with that fixture's
+   `workspace_prefix` are rewritten to `../<as>` so they resolve next to the
+   skill. Upstream files are not modified.
 3b. Optionally vendors the skill under a different local catalog name (the
    `as:` field on a skill entry). Federated skills follow a
    `<projectrepo>-<skill>` naming convention in this catalog (e.g. the
@@ -143,8 +149,10 @@
 MARKETPLACE_DESCRIPTION_MAX = 320
 
 REPO_RE = re.compile(r"^[A-Za-z0-9._-]+/[A-Za-z0-9._-]+$")
-SOURCE_KEYS = {"repo", "license", "branch", "skills"}
+SOURCE_KEYS = {"repo", "license", "branch", "skills", "fixtures"}
 SKILL_KEYS = {"path", "as", "marketplace_description"}
+FIXTURE_KEYS = {"path", "as", "workspace_prefix", "files"}
+FIXTURE_MARKER = ".fixture.json"
 # Keys that would pin a source to something other than a branch. They are
 # rejected with a pointed message rather than ignored, so an entry that tries
 # to track a tag or commit fails loudly instead of silently tracking `main`.
@@ -170,11 +178,28 @@ def dest_name(self) -> str:
         return self.local_name or self.folder
 
 
+@dataclass
+class FixtureSpec:
+    """A non-skill directory vendored once and shared by a source's skills.
+
+    `workspace_prefix` is the path those skills use in their own repo
+    (for example ``../../sample``). Vendored datasets are rewritten so the
+    same value resolves to ``../<local_name>`` beside the skill. `files` are
+    the paths under `path` that get vendored; nothing else is copied.
+    """
+
+    path: str
+    local_name: str
+    workspace_prefix: str
+    files: list[str] = field(default_factory=list)
+
+
 @dataclass
 class Source:
     repo: str
     license: str
     skills: list[SkillSpec] = field(default_factory=list)
+    fixtures: list[FixtureSpec] = field(default_factory=list)
     # As declared: a branch name, or a pattern such as `release/*`.
     branch: str = branches.DEFAULT_BRANCH
 
@@ -295,17 +320,107 @@ def parse_federation(catalog: Path) -> list[Source]:
                 )
             )
 
+        fixtures = _parse_fixtures(raw.get("fixtures", []), f"{where}.fixtures")
         sources.append(
             Source(
                 repo=repo,
                 license=raw.get("license") or "UNKNOWN",
                 skills=skills,
+                fixtures=fixtures,
                 branch=branch,
             )
         )
+    _reject_name_collisions(sources)
     return sources
 
 
+def _parse_repo_path(path: object, where: str) -> str:
+    """A repo-relative directory path: no scheme, no ``..``, no leading slash."""
+    if not isinstance(path, str) or not path.strip("/"):
+        raise ValueError(
+            f"{where} must be a directory path inside the repo, got {path!r}."
+        )
+    normalized = path.replace("\\", "/").strip("/")
+    if ".." in normalized.split("/"):
+        raise ValueError(f"{where} must not contain '..': {path!r}")
+    return normalized
+
+
+def _parse_fixtures(raw: object, where: str) -> list[FixtureSpec]:
+    if raw is None:
+        return []
+    if not isinstance(raw, list):
+        raise ValueError(f"{where} must be an array.")
+    fixtures: list[FixtureSpec] = []
+    for idx, item in enumerate(raw):
+        item_where = f"{where}[{idx}]"
+        if not isinstance(item, dict):
+            raise ValueError(f"{item_where} must be an object.")
+        unknown = set(item) - FIXTURE_KEYS
+        if unknown:
+            raise ValueError(
+                f"{item_where} has unknown key(s) {sorted(unknown)}; expected "
+                f"{sorted(FIXTURE_KEYS)}."
+            )
+        local_name = item.get("as")
+        if not isinstance(local_name, str) or not local_name or "/" in local_name:
+            raise ValueError(
+                f"{item_where}.as must be a plain local folder name, got "
+                f"{local_name!r}."
+            )
+        if local_name.startswith("."):
+            raise ValueError(f"{item_where}.as must not be hidden: {local_name!r}")
+        prefix = item.get("workspace_prefix")
+        if not isinstance(prefix, str) or not prefix.strip("/"):
+            raise ValueError(
+                f"{item_where}.workspace_prefix must be the relative workspace "
+                f"path used upstream, got {prefix!r}."
+            )
+        prefix = prefix.replace("\\", "/").strip("/")
+        if prefix.startswith("/") or "://" in prefix:
+            raise ValueError(
+                f"{item_where}.workspace_prefix must be a relative path, got "
+                f"{prefix!r}."
+            )
+        files = item.get("files")
+        if not isinstance(files, list) or not files:
+            raise ValueError(
+                f"{item_where}.files must list the fixture files to vendor."
+            )
+        fixtures.append(
+            FixtureSpec(
+                path=_parse_repo_path(item.get("path"), f"{item_where}.path"),
+                local_name=local_name,
+                workspace_prefix=prefix,
+                files=[
+                    _parse_repo_path(rel, f"{item_where}.files[{i}]")
+                    for i, rel in enumerate(files)
+                ],
+            )
+        )
+    return fixtures
+
+
+def _reject_name_collisions(sources: list[Source]) -> None:
+    """Skill folders and fixture folders share `skills/`, so names must be unique."""
+    seen: set[str] = set()
+    for source in sources:
+        for spec in source.skills:
+            if spec.dest_name in seen:
+                raise ValueError(
+                    f"Skill name collision: {spec.dest_name!r} is declared "
+                    "more than once."
+                )
+            seen.add(spec.dest_name)
+        for fixture in source.fixtures:
+            if fixture.local_name in seen:
+                raise ValueError(
+                    f"Fixture name collision: {fixture.local_name!r} is also "
+                    "used as a skill or fixture folder."
+                )
+            seen.add(fixture.local_name)
+
+
 def run(cmd: list[str], cwd: Path | None = None) -> str:
     """Run a command, raise on failure, return stdout."""
     result = subprocess.run(
@@ -746,7 +861,11 @@ def import_source(
             log.append(f"[{source.name}] {source.branch} resolved to {branch}")
         log.append(f"[{source.name}] cloning {source.repo}@{branch}")
         commit = shallow_clone(
-            source.repo, branch, [spec.path for spec in source.skills], tmp_path
+            source.repo,
+            branch,
+            [spec.path for spec in source.skills]
+            + [fixture.path for fixture in source.fixtures],
+            tmp_path,
         )
         log.append(f"[{source.name}] resolved to commit {commit}")
         repo_files = list_repo_files(tmp_path, commit)
@@ -873,9 +992,118 @@ def import_source(
                     omitted=omitted,
                 )
             )
+        if source.fixtures:
+            sync_fixtures(source, tmp_path, branch, commit, log)
     return results
 
 
+def rewrite_fixture_workspaces(skill_dir: Path, fixture: FixtureSpec) -> bool:
+    """Point vendored workspace paths at the catalog copy of a fixture.
+
+    Upstream keeps ``workspace_prefix`` (``../../sample`` inside the product
+    tree). The catalog copy lives beside the skill, so the same cases use
+    ``../<fixture>``. A second run finds nothing to replace.
+    """
+    prefix = fixture.workspace_prefix
+    replacement = f"../{fixture.local_name}"
+    changed = False
+    for name in ("evals/evals.json", "evals/extended_evals.json"):
+        path = skill_dir / name
+        if not path.is_file():
+            continue
+        original = path.read_text(encoding="utf-8")
+        updated = original.replace(f'"{prefix}', f'"{replacement}')
+        if updated != original:
+            path.write_text(updated, encoding="utf-8")
+            changed = True
+    return changed
+
+
+def fixture_hash(root: Path, files: Iterable[str]) -> str:
+    """Hash the named files under `root`, in the same form as `content_hash`."""
+    digest = hashlib.sha256()
+    for rel in sorted(files):
+        digest.update(rel.encode("utf-8"))
+        digest.update(b"\0")
+        digest.update(hashlib.sha256((root / rel).read_bytes()).digest())
+    return digest.hexdigest()
+
+
+def sync_fixtures(
+    source: Source,
+    clone_dir: Path,
+    branch: str,
+    commit: str,
+    log: list[str],
+    skills_dir: Path = SKILLS_DIR,
+) -> None:
+    """Copy each declared fixture under `skills/` and retarget workspace paths.
+
+    The fixture is not a skill: its marker is `.fixture.json`, so the
+    undeclared-skill report does not treat the directory as a vendored skill.
+    Only the fixture's declared `files` are copied, so the vendored copy holds
+    what the eval workspaces use rather than the whole upstream tree.
+    """
+    for fixture in source.fixtures:
+        src = clone_dir / fixture.path
+        missing = [rel for rel in fixture.files if not (src / rel).is_file()]
+        if missing:
+            raise FileNotFoundError(
+                f"Fixture file(s) {missing} not found under {fixture.path!r} in "
+                f"{source.repo}@{branch}."
+            )
+        dest = skills_dir / fixture.local_name
+        upstream_hash = fixture_hash(src, fixture.files)
+        marker_path = dest / FIXTURE_MARKER
+        previous = {}
+        if marker_path.is_file():
+            try:
+                loaded = json.loads(marker_path.read_text(encoding="utf-8"))
+                previous = loaded if isinstance(loaded, dict) else {}
+            except json.JSONDecodeError:
+                previous = {}
+        same = (
+            previous.get("content_hash") == upstream_hash
+            and previous.get("repo") == source.repo
+            and previous.get("path") == fixture.path
+            and previous.get("ref") == source.branch
+        )
+        if same:
+            log.append(
+                f"[{source.name}] skills/{fixture.local_name} fixture is up to date"
+            )
+        else:
+            if dest.exists():
+                shutil.rmtree(dest)
+            for rel in fixture.files:
+                (dest / rel).parent.mkdir(parents=True, exist_ok=True)
+                shutil.copy2(src / rel, dest / rel)
+            marker = {
+                "repo": source.repo,
+                "ref": source.branch,
+                "commit": commit,
+                "path": fixture.path,
+                "content_hash": upstream_hash,
+                "imported_at": datetime.now(timezone.utc).strftime("%Y-%m-%dT%H:%M:%SZ"),
+            }
+            if branch != source.branch:
+                marker["resolved_ref"] = branch
+            (dest / FIXTURE_MARKER).write_text(
+                json.dumps(marker, indent=2) + "\n", encoding="utf-8"
+            )
+            log.append(
+                f"[{source.name}] vendoring fixture {fixture.path} -> "
+                f"skills/{fixture.local_name}"
+            )
+        for spec in source.skills:
+            skill_dir = skills_dir / spec.dest_name
+            if rewrite_fixture_workspaces(skill_dir, fixture):
+                log.append(
+                    f"[{source.name}] skills/{spec.dest_name}: workspace "
+                    f"{fixture.workspace_prefix} -> ../{fixture.local_name}"
+                )
+
+
 def report_undeclared(
     declared: set[str],
     existing: dict[str, dict],
@@ -1049,6 +1277,13 @@ def main(argv: list[str] | None = None) -> int:
                     f"{source.repo}@{source.branch} {spec.path} "
                     f"-> skills/{spec.dest_name}"
                 )
+            for fixture in source.fixtures:
+                print(
+                    f"{source.repo}@{source.branch} {fixture.path} "
+                    f"-> skills/{fixture.local_name} "
+                    f"({len(fixture.files)} file(s), workspace "
+                    f"{fixture.workspace_prefix})"
+                )
         print("")
         print(f"{args.catalog} is valid: {len(seen)} skill(s) declared.")
         return 0
diff --git a/.github/scripts/test_federate_skills.py b/.github/scripts/test_federate_skills.py
index f3d374e6..00e33f11 100644
--- a/.github/scripts/test_federate_skills.py
+++ b/.github/scripts/test_federate_skills.py
@@ -147,6 +147,129 @@ def test_rejects_malformed_entries(self):
                 with self.assertRaises(ValueError):
                     parse(payload)
 
+    def test_parses_a_shared_fixture(self):
+        source = parse(
+            one_source(
+                fixtures=[
+                    {
+                        "path": "projects/rocprofiler-compute/sample",
+                        "as": "rocprof-compute-sample",
+                        "workspace_prefix": "../../sample",
+                        "files": ["vcopy.cpp", "wmma_gemm_fp16/Makefile"],
+                    }
+                ]
+            )
+        )[0]
+        self.assertEqual(len(source.fixtures), 1)
+        fixture = source.fixtures[0]
+        self.assertEqual(fixture.path, "projects/rocprofiler-compute/sample")
+        self.assertEqual(fixture.local_name, "rocprof-compute-sample")
+        self.assertEqual(fixture.workspace_prefix, "../../sample")
+        self.assertEqual(fixture.files, ["vcopy.cpp", "wmma_gemm_fp16/Makefile"])
+
+    def test_rejects_a_fixture_without_files_or_with_dotdot(self):
+        base = {
+            "path": "projects/sample",
+            "as": "rocprof-compute-sample",
+            "workspace_prefix": "../../sample",
+        }
+        for files in (None, [], ["../secret.txt"]):
+            with self.subTest(files=files):
+                entry = dict(base) if files is None else {**base, "files": files}
+                with self.assertRaises(ValueError):
+                    parse(one_source(fixtures=[entry]))
+
+    def test_rejects_a_fixture_that_collides_with_a_skill(self):
+        with self.assertRaises(ValueError):
+            parse(
+                one_source(
+                    skills=[{"path": "skills/sample", "as": "rocprof-compute-sample"}],
+                    fixtures=[
+                        {
+                            "path": "projects/sample",
+                            "as": "rocprof-compute-sample",
+                            "workspace_prefix": "../../sample",
+                            "files": ["vcopy.cpp"],
+                        }
+                    ],
+                )
+            )
+
+    def test_rewrites_workspace_prefixes_once(self):
+        fixture = fed.FixtureSpec(
+            path="projects/sample",
+            local_name="rocprof-compute-sample",
+            workspace_prefix="../../sample",
+        )
+        with tempfile.TemporaryDirectory() as tmp:
+            skill = write_skill(
+                Path(tmp) / "skill",
+                {
+                    "evals/evals.json": (
+                        '{ "workspace": "../../sample/wmma_gemm_fp16" }\n'
+                    )
+                },
+            )
+            self.assertTrue(fed.rewrite_fixture_workspaces(skill, fixture))
+            text = (skill / "evals" / "evals.json").read_text(encoding="utf-8")
+            self.assertIn("../rocprof-compute-sample/wmma_gemm_fp16", text)
+            self.assertNotIn("../../sample", text)
+            self.assertFalse(fed.rewrite_fixture_workspaces(skill, fixture))
+
+    def test_sync_fixtures_copies_listed_files_and_retargets_evals(self):
+        source = fed.Source(
+            repo="ROCm/rocm-systems",
+            license="MIT",
+            branch="develop",
+            skills=[fed.SkillSpec(path="skills/demo", local_name="demo")],
+            fixtures=[
+                fed.FixtureSpec(
+                    path="projects/sample",
+                    local_name="rocprof-compute-sample",
+                    workspace_prefix="../../sample",
+                    files=["vcopy.cpp", "wmma/Makefile"],
+                )
+            ],
+        )
+        with tempfile.TemporaryDirectory() as tmp:
+            root = Path(tmp)
+            write_skill(
+                root / "clone" / "projects" / "sample",
+                {
+                    "vcopy.cpp": "int main(){return 0;}\n",
+                    "wmma/Makefile": "all:\n",
+                    "mega_kernel/main.cpp": "int main(){return 1;}\n",
+                },
+            )
+            write_skill(
+                root / "skills" / "demo",
+                {"evals/evals.json": '{ "workspace": "../../sample" }\n'},
+            )
+            log: list[str] = []
+            fed.sync_fixtures(
+                source, root / "clone", "develop", "abc123", log, root / "skills"
+            )
+            copied = root / "skills" / "rocprof-compute-sample" / "vcopy.cpp"
+            self.assertEqual(copied.read_text(encoding="utf-8"), "int main(){return 0;}\n")
+            vendored = root / "skills" / "rocprof-compute-sample"
+            self.assertTrue((vendored / "wmma" / "Makefile").is_file())
+            self.assertFalse((vendored / "mega_kernel").exists())
+            marker = json.loads(
+                (root / "skills" / "rocprof-compute-sample" / ".fixture.json").read_text(
+                    encoding="utf-8"
+                )
+            )
+            self.assertEqual(marker["commit"], "abc123")
+            self.assertEqual(marker["path"], "projects/sample")
+            evals = (root / "skills" / "demo" / "evals" / "evals.json").read_text(
+                encoding="utf-8"
+            )
+            self.assertIn("../rocprof-compute-sample", evals)
+            fed.sync_fixtures(
+                source, root / "clone", "develop", "abc123", log, root / "skills"
+            )
+            self.assertTrue(any("up to date" in line for line in log))
+
     def test_accepts_windows_style_separators_in_paths(self):
         skills = parse(
             one_source(skills=[{"path": r"TraceLens\Agent\skills\orchestrator"}])
diff --git a/CONTRIBUTING.md b/CONTRIBUTING.md
index 7d9250ac..8e6c1cde 100644
--- a/CONTRIBUTING.md
+++ b/CONTRIBUTING.md
@@ -101,6 +101,7 @@ one repo can federate as many skills as it likes from wherever they live:
 | `branch` | Optional. The branch to track, or a release pattern such as `release/*`. Defaults to `main` |
 | `skills[].path` | Path of the skill folder inside your repo, from the repo root |
 | `skills[].as` | Optional local catalog name; use it to namespace as `<project>-<skill>` so names stay unique |
+| `fixtures[]` | Optional directory that is not a skill. Only the paths in `files` are copied, once, to `skills/<as>/`. `workspace_prefix` is the relative `workspace` value in the upstream datasets (`../../sample`). Vendored datasets are rewritten to `../<as>` |
 
 There is no tag or commit field. A source always follows a branch.
 
diff --git a/skills/rocprof-compute-kernel-bottleneck/.federated.json b/skills/rocprof-compute-kernel-bottleneck/.federated.json
new file mode 100644
index 00000000..3b09dff4
--- /dev/null
+++ b/skills/rocprof-compute-kernel-bottleneck/.federated.json
@@ -0,0 +1,10 @@
+{
+  "source": "rocm-rocm-systems",
+  "repo": "ROCm/rocm-systems",
+  "ref": "develop",
+  "commit": "8cce7054f17a77272e1ba73ef78f4f77cf056956",
+  "path": "projects/rocprofiler-compute/skills/rocprof-compute-kernel-bottleneck",
+  "license": "MIT",
+  "content_hash": "fc6c0e7c031c170e9f2cff67973a40ec8c9ccfecbd513f057c79121fd9e2a2c2",
+  "imported_at": "2026-10-28T01:19:39Z"
+}
diff --git a/skills/rocprof-compute-kernel-bottleneck/SKILL.md b/skills/rocprof-compute-kernel-bottleneck/SKILL.md
new file mode 100644
index 00000000..35ea364b
--- /dev/null
+++ b/skills/rocprof-compute-kernel-bottleneck/SKILL.md
@@ -0,0 +1,244 @@
+---
+name: rocprof-compute-kernel-bottleneck
+description: Profiles an AMD GPU application with rocprof-compute and finds which kernel is slow and why, using performance metrics and PC sampling. Collects counters and selects kernels and dispatches, then analyzes GPU memory throughput, compute efficiency, wavefront occupancy and wavefront limits, scheduler stalls, etc. Use when the user asks to profile, benchmark, or speed up a HIP/ROCm kernel or application, or asks where a kernel's bottleneck is. Not for CUDA tools, Windows even when the user names rocprof-compute, system-wide CPU/MPI tracing, or writing kernel source.
+---
+
+# Profile an AMD GPU application and find the kernel bottleneck
+
+Start here for any "profile this" or "why is this kernel slow" request. This
+skill collects counters, narrows a workload down to one kernel, and identifies
+what limits it. Hand off to a focused skill once the limit is clear.
+
+Run `rocprof-compute profile --help` and `rocprof-compute analyze --help`
+before choosing flags. Options change between releases; do not guess one.
+
+## Two sources of performance data
+
+rocprof-compute collects from two independent sources, and they answer
+different questions. Choose by what the user is asking, not by difficulty.
+
+| Source | Answers | Skill |
+|---|---|---|
+| Perfmon counters | How the architecture behaves: bandwidth, cache, occupancy, pipeline utilization | this skill, then `rocprof-compute-speed-of-light`, `rocprof-compute-memory`, `rocprof-compute-roofline` |
+| PC sampling | How the source code behaves: which instruction or line is hot and why it stalls | `rocprof-compute-pc-sampling` |
+
+Counters are the default starting point because they cover the whole kernel
+cheaply. Go straight to `rocprof-compute-pc-sampling` when the user asks about instructions,
+source lines, or stall reasons, and use both when an architectural limit needs
+to be traced back to the code that causes it. A request to run `rocprofv3`
+itself is neither this skill nor `rocprof-compute-pc-sampling`.
+
+## 1. Check the environment
+
+In this order, because each command depends on the one before it:
+
+```bash
+amd-smi static
+amd-smi list
+rocminfo
+rocprof-compute --specs
+```
+
+`amd-smi` and `rocminfo` report the GPU and ROCm install.
+`rocprof-compute --specs` reports what the profiler itself detected, including
+`ROCm Version`, `GPU Model`, and `GPU Arch`. It is built on the first two, so
+it cannot succeed when they fail.
+
+Supported accelerators are listed in
+[compatible-accelerators.rst](https://github.com/ROCm/rocm-systems/blob/8cce7054f17a77272e1ba73ef78f4f77cf056956/projects/rocprofiler-compute/docs/reference/compatible-accelerators.rst).
+Not every feature works on every architecture. Check there before promising a
+result, and let the tool's own error message stand when a feature is
+unavailable.
+
+## 2. Collect what you need before profiling
+
+Ask once:
+
+1. Launch command and arguments
+2. A short workload name, used as `--name`
+3. Which kernels matter, if not all
+
+Profiling replays the application once per counter pass, so the command must be
+safe to run repeatedly and short enough to run several times. Ask for a reduced
+iteration count or problem size rather than profiling a long run. Do not invent
+a timeout; profile mode has no such option.
+
+## 3. Profile
+
+```bash
+rocprof-compute profile --name <workload_name> -- <application> <args>
+```
+
+Output lands in `./workloads/<workload_name>/<gpu_model>/`, or `<rank>/` under
+MPI. Change it with `--output-directory`, which replaces `--name` for building
+the path. Re-profiling into a non-empty directory needs `--overwrite`; prefer a
+fresh directory.
+
+Which GPU the application runs on is in
+[Selecting a GPU](https://github.com/ROCm/rocm-systems/blob/8cce7054f17a77272e1ba73ef78f4f77cf056956/projects/rocprofiler-compute/docs/how-to/profile/mode.rst). Under Slurm, derive
+`HIP_VISIBLE_DEVICES` from `ROCR_VISIBLE_DEVICES`, not from `SLURM_STEP_GPUS`;
+the indices differ.
+
+See [profile mode](https://github.com/ROCm/rocm-systems/blob/8cce7054f17a77272e1ba73ef78f4f77cf056956/projects/rocprofiler-compute/docs/how-to/profile/mode.rst) for output layout,
+filtering, and the full option list. To profile a process that is already
+running, see
+[live attach and detach](https://github.com/ROCm/rocm-systems/blob/8cce7054f17a77272e1ba73ef78f4f77cf056956/projects/rocprofiler-compute/docs/how-to/live_attach_detach.rst).
+
+### Collect less when a full replay is too expensive
+
+A single-pass, topic-focused collection:
+
+```bash
+rocprof-compute profile --list-sets
+rocprof-compute profile --name <name> --set <set_name> -- <application>
+```
+
+Or collect only named blocks:
+
+```bash
+rocprof-compute profile --name <name> -b sol wavefront -- <application>
+```
+
+Iteration multiplexing spreads counters across kernel launches so a single pass
+collects more:
+
+```bash
+rocprof-compute profile --name <name> --iteration-multiplexing kernel -- <application>
+```
+
+Whenever you collect a subset, tell the user which analysis panels will be
+empty and that answering those questions needs another profile.
+
+## 4. Find the hot kernel
+
+```bash
+rocprof-compute analyze --path ./workloads/<name>/<gpu_model> --list-stats
+```
+
+`--path` must point at the directory holding `profiling_config.yaml`. Analyze
+does not search for it in nearby directories. If the path you have does not
+contain that file, look one level down for the GPU model or rank directory.
+
+`--list-stats` prints `Detected Kernels (sorted descending by duration)` and a
+`Dispatch list`. The kernel index is what `-k` takes. Dispatch ids are 1-based
+and go to `-d`.
+
+`--list-stats` cannot be used with `--output-format csv` or `db`.
+
+Work on one kernel at a time:
+
+```bash
+rocprof-compute analyze --path ./workloads/<name>/<gpu_model> -k 0
+```
+
+`-k` at profile time is different: it filters by kernel name and is used as a
+regex, so `-k gemm` matches every kernel whose name contains `gemm`.
+
+## 5. Narrow to a reason
+
+Get the overview, then open exactly one detailed block:
+
+```bash
+rocprof-compute analyze --path ./workloads/<name>/<gpu_model> -k 0 -b sol
+```
+
+Read the `Percent of Peak` column. The report prints the peak for each metric,
+so compare against that rather than against a remembered threshold, and never
+compute a peak yourself.
+
+Resolve block names with aliases, which are stable across releases, instead of
+numeric ids, which are architecture-specific:
+
+```bash
+rocprof-compute --list-blocks <arch>
+rocprof-compute analyze --path ./workloads/<name>/<gpu_model> --list-available-metrics
+```
+
+Where to go next:
+
+| Symptom in Speed-of-Light | Next step |
+|---|---|
+| Memory throughput near peak, low compute | the `rocprof-compute-memory` skill |
+| Compute near peak, or unclear compute vs memory | the `rocprof-compute-roofline` skill |
+| Everything far below peak | `-b wavefront` for occupancy, then `-b spi` for launch limits, both below |
+| Question is about instructions, source lines, or stall reasons | the `rocprof-compute-pc-sampling` skill |
+| Kernels come from PyTorch | the `rocprof-compute-torch-trace` skill |
+
+### Occupancy and wavefront limits
+
+```bash
+rocprof-compute analyze --path ./workloads/<name>/<gpu_model> -k 0 -b wavefront
+```
+
+Compare active waves per CU against the theoretical occupancy in the same
+table. When theoretical occupancy is itself low, the limit is VGPR, SGPR, or
+LDS allocation. Scratch traffic points at register spilling. Low occupancy is
+not a defect when the kernel already saturates its limiting resource. Worked
+examples are in
+[occupancy-limiters-example.rst](https://github.com/ROCm/rocm-systems/blob/8cce7054f17a77272e1ba73ef78f4f77cf056956/projects/rocprofiler-compute/docs/tutorial/includes/occupancy-limiters-example.rst).
+
+### Scheduler and pipeline limits
+
+```bash
+# Workgroup launch rate and resource allocation limits
+rocprof-compute analyze --path ./workloads/<name>/<gpu_model> -k 0 -b spi
+
+# Pipeline utilization and instruction mix
+rocprof-compute analyze --path ./workloads/<name>/<gpu_model> -k 0 -b cu_pipe
+```
+
+There is no single stall panel. Stall counters sit in the block that owns the
+hardware: LDS stalls in `lds`, cache stalls in `vl1d` and `l2`, fabric stalls
+in `l2_per_channel`. Open the block the Speed-of-Light result points at, and
+read what each counter means in the
+[performance model](https://github.com/ROCm/rocm-systems/blob/8cce7054f17a77272e1ba73ef78f4f77cf056956/projects/rocprofiler-compute/docs/conceptual/performance-model.rst). Ground every
+recommendation in a number you read from the report.
+
+## 6. Normalization
+
+`-n` / `--normal-unit` applies to every metric in the report at once:
+`per_kernel` (default), `per_wave`, `per_cycle`, `per_second`. Keep the default
+for a mixed report. Change it only when the user asks about one metric family
+and the unit still makes sense for it. Bandwidth per cycle does not.
+
+## 7. Compare two runs
+
+Pass `--path` twice to put two workloads side by side. Use this for:
+
+1. Measuring a software optimization against its baseline on the same GPU.
+2. Checking that the same code performs consistently across nodes with the
+   same GPU architecture.
+3. Comparing the same code across different GPU architectures.
+
+```bash
+rocprof-compute analyze \
+    --path ./workloads/baseline/<gpu_model> \
+    --path ./workloads/optimized/<gpu_model>
+```
+
+Cases 1 and 2 expect the numbers to match except where the change intended
+otherwise, so an unexplained difference is the finding. Case 3 compares
+different hardware, so compare percent of peak rather than absolute values.
+
+Never present a comparison across different GPUs, partition modes, or clock
+settings without stating that those differ.
+
+## 8. Save a report
+
+Terminal output is the default. Only choose a format when the user asks to save
+or export. `--output-format` takes `txt`, `csv`, or `db`, each of which
+disables terminal output, and `--output-name` sets the file name. See
+[analysis output format](https://github.com/ROCm/rocm-systems/blob/8cce7054f17a77272e1ba73ef78f4f77cf056956/projects/rocprofiler-compute/docs/how-to/analyze/cli.rst).
+
+## 9. When something looks wrong
+
+| Symptom | Likely cause |
+|---|---|
+| No workload data | `--path` is not the directory with `profiling_config.yaml` |
+| Empty or missing panels | those counters were not collected; re-profile |
+| All metrics zero | kernel too short, or no dispatches were captured |
+| No kernel ids | the application launched no GPU kernels |
+| Roofline missing | not collected, or unsupported on this architecture |
+
+Stop and report permission, firmware, and ROCm version failures instead of
+working around them. More in the [FAQ](https://github.com/ROCm/rocm-systems/blob/8cce7054f17a77272e1ba73ef78f4f77cf056956/projects/rocprofiler-compute/docs/reference/faq.rst).
diff --git a/skills/rocprof-compute-kernel-bottleneck/evals/evals.json b/skills/rocprof-compute-kernel-bottleneck/evals/evals.json
new file mode 100644
index 00000000..a13db25f
--- /dev/null
+++ b/skills/rocprof-compute-kernel-bottleneck/evals/evals.json
@@ -0,0 +1,105 @@
+{
+  "evaluations": [
+    {
+      "id": "kb-profile-hip-application",
+      "skill_should_trigger": true,
+      "prompt": "profile my HIP application with rocprof-compute and tell me where the time goes\n\nThe application is ./vcopy.cpp.",
+      "workspace": "../rocprof-compute-sample",
+      "expected_behavior": [
+        "asks for the launch command and a workload name before profiling",
+        "runs rocprof-compute profile --name <name> -- <application>",
+        "runs analyze with --list-stats to find the hottest kernel"
+      ],
+      "unexpected_behavior": [
+        "invents a profiling flag instead of checking --help"
+      ],
+      "logs_contain": ["rocprof-compute profile", "--list-stats"]
+    },
+    {
+      "id": "kb-which-kernel-is-slow",
+      "skill_should_trigger": true,
+      "prompt": "which kernel is the hotspot in ./workloads/my_app and why is it slow?\n\nThe application is ./vmem.hip.",
+      "workspace": "../rocprof-compute-sample",
+      "expected_behavior": [
+        "points --path at the directory containing profiling_config.yaml rather than a parent directory",
+        "uses --list-stats first, then narrows to one kernel id with -k",
+        "reads the System Speed-of-Light panel before opening a detailed block"
+      ],
+      "logs_contain": ["--list-stats", "-k"]
+    },
+    {
+      "id": "kb-occupancy-limited",
+      "skill_should_trigger": true,
+      "prompt": "is occupancy limiting this rocprof-compute kernel?\n\nThe application is ./occupancy.hip.",
+      "workspace": "../rocprof-compute-sample",
+      "expected_behavior": [
+        "opens the wavefront block for the selected kernel",
+        "compares achieved occupancy against the theoretical occupancy printed in the report",
+        "ties any VGPR, SGPR, or LDS advice to numbers read from the report"
+      ],
+      "unexpected_behavior": [
+        "calls low occupancy a defect without checking whether the kernel already saturates another resource"
+      ],
+      "logs_contain": ["wavefront"]
+    },
+    {
+      "id": "kb-expensive-replay-fallback",
+      "skill_should_trigger": true,
+      "prompt": "a full rocprof-compute replay of this workload is too expensive, collect something cheaper\n\nThe application is ./vmem.hip.",
+      "workspace": "../rocprof-compute-sample",
+      "expected_behavior": [
+        "offers a single-pass option such as --set, a narrowed -b, or --iteration-multiplexing",
+        "runs --list-sets before selecting a set",
+        "warns which analysis panels will be empty afterwards"
+      ],
+      "logs_contain": ["--list-sets"]
+    },
+    {
+      "id": "kb-unsafe-to-replay",
+      "skill_should_trigger": true,
+      "prompt": "profile this one-shot ROCm job that writes to the production database",
+      "expected_behavior": [
+        "warns that counter profiling replays the application multiple times",
+        "does not start profiling until replay safety is settled",
+        "offers a single-pass collection as an alternative"
+      ],
+      "unexpected_behavior": [
+        "runs a multi-pass profile of a workload the user described as unsafe to repeat"
+      ]
+    },
+    {
+      "id": "kb-kernel-name-filter",
+      "skill_should_trigger": true,
+      "prompt": "profile only the gemm kernels in my ROCm application\n\nThe application is ./wmma_gemm_fp16.hip.",
+      "workspace": "../rocprof-compute-sample/wmma_gemm_fp16",
+      "expected_behavior": [
+        "uses rocprof-compute profile -k to filter by kernel name",
+        "treats the profile-time -k value as a regex matched against kernel names"
+      ],
+      "logs_contain": ["-k"]
+    },
+    {
+      "id": "kb-cuda-application",
+      "skill_should_trigger": false,
+      "note": "Profiling request, but CUDA tooling is out of scope.",
+      "prompt": "profile my CUDA application with nsys"
+    },
+    {
+      "id": "kb-windows-request",
+      "skill_should_trigger": false,
+      "note": "rocprof-compute is named, but the platform is unsupported.",
+      "prompt": "profile my application on Windows with rocprof-compute"
+    },
+    {
+      "id": "kb-install-question",
+      "skill_should_trigger": false,
+      "prompt": "how do I install ROCm on Ubuntu?"
+    },
+    {
+      "id": "kb-write-a-kernel",
+      "skill_should_trigger": false,
+      "note": "GPU performance adjacent, but this skill does not author kernel source.",
+      "prompt": "write me a faster HIP matrix multiply kernel"
+    }
+  ]
+}
diff --git a/skills/rocprof-compute-kernel-bottleneck/evals/machine.yml b/skills/rocprof-compute-kernel-bottleneck/evals/machine.yml
new file mode 100644
index 00000000..4d7efd45
--- /dev/null
+++ b/skills/rocprof-compute-kernel-bottleneck/evals/machine.yml
@@ -0,0 +1,2 @@
+os: [Linux]
+labels: [mi300x, gpu, rocm]
diff --git a/skills/rocprof-compute-kernel-bottleneck/skill-card.md b/skills/rocprof-compute-kernel-bottleneck/skill-card.md
new file mode 100644
index 00000000..0b0d344d
--- /dev/null
+++ b/skills/rocprof-compute-kernel-bottleneck/skill-card.md
@@ -0,0 +1,13 @@
+# Skill Card
+
+## Description
+
+Profiles an application running on an AMD GPU with rocprof-compute and identifies which kernel is the bottleneck and why.
+
+## Owner
+
+rocprofiler-compute, a project in [rocm-systems](https://github.com/ROCm/rocm-systems/)
+
+## License
+
+MIT
diff --git a/skills/rocprof-compute-memory/.federated.json b/skills/rocprof-compute-memory/.federated.json
new file mode 100644
index 00000000..8fbf5fcf
--- /dev/null
+++ b/skills/rocprof-compute-memory/.federated.json
@@ -0,0 +1,10 @@
+{
+  "source": "rocm-rocm-systems",
+  "repo": "ROCm/rocm-systems",
+  "ref": "develop",
+  "commit": "8cce7054f17a77272e1ba73ef78f4f77cf056956",
+  "path": "projects/rocprofiler-compute/skills/rocprof-compute-memory",
+  "license": "MIT",
+  "content_hash": "78037d87af423a68edab4d251b787cc07763467f570214c657bade7fecc01c74",
+  "imported_at": "2026-10-28T01:19:39Z"
+}
diff --git a/skills/rocprof-compute-memory/SKILL.md b/skills/rocprof-compute-memory/SKILL.md
new file mode 100644
index 00000000..1f88ca19
--- /dev/null
+++ b/skills/rocprof-compute-memory/SKILL.md
@@ -0,0 +1,117 @@
+---
+name: rocprof-compute-memory
+description: Analyzes the AMD GPU memory hierarchy with rocprof-compute, covering the Memory Chart, vL1D and L2 caches, LDS bank conflicts, Infinity Fabric traffic, and the experimental gfx950 memory bandwidth analysis. Use when the user asks about memory bandwidth, cache hit rates, coalescing, LDS conflicts, data locality, HBM traffic, or says a kernel is memory-bound. Not for CUDA tools, Nsight, Windows, host-memory leaks, GPU out-of-memory allocation failures, or system-wide tracing.
+---
+
+# Analyze the AMD GPU memory hierarchy
+
+Use this after Speed-of-Light points at memory, or when the user asks about
+bandwidth, caches, or data movement directly.
+
+Run `rocprof-compute analyze --help` before choosing flags.
+
+## 1. Start with the Memory Chart
+
+```bash
+rocprof-compute analyze --path ./workloads/<name>/<gpu_model> --list-stats
+rocprof-compute analyze --path ./workloads/<name>/<gpu_model> -k <kernel_id> -b memchart
+```
+
+The Memory Chart shows the whole hierarchy at once: requests issued, what each
+cache level served, and what reached HBM. Read it before opening any single
+cache level, because it tells you which level to open.
+
+Always analyze one kernel at a time with `-k`.
+
+## 2. Open the level the chart points at
+
+```bash
+# Vector L1 data cache
+rocprof-compute analyze --path ./workloads/<name>/<gpu_model> -k <kernel_id> -b vl1d
+
+# L2 cache
+rocprof-compute analyze --path ./workloads/<name>/<gpu_model> -k <kernel_id> -b l2
+
+# L2 per channel, including Infinity Fabric read and write stalls
+rocprof-compute analyze --path ./workloads/<name>/<gpu_model> -k <kernel_id> -b l2_per_channel
+
+# Local Data Share, including bank conflicts and unaligned stalls
+rocprof-compute analyze --path ./workloads/<name>/<gpu_model> -k <kernel_id> -b lds
+
+# Address processing and data return, for coalescing behaviour
+rocprof-compute analyze --path ./workloads/<name>/<gpu_model> -k <kernel_id> -b tatd
+
+# Scalar L1 data cache
+rocprof-compute analyze --path ./workloads/<name>/<gpu_model> -k <kernel_id> -b sl1d
+```
+
+Aliases are stable; numeric block ids are architecture-specific. List what this
+workload actually has with `--list-available-metrics`, and the aliases for an
+architecture with `rocprof-compute --list-blocks <arch>`.
+
+## 3. Interpret
+
+Each level has its own reference, with the counters, the formulas, and what
+they mean:
+
+- [Vector L1 cache](https://github.com/ROCm/rocm-systems/blob/8cce7054f17a77272e1ba73ef78f4f77cf056956/projects/rocprofiler-compute/docs/conceptual/cdna/vector-l1-cache.rst)
+- [L2 cache](https://github.com/ROCm/rocm-systems/blob/8cce7054f17a77272e1ba73ef78f4f77cf056956/projects/rocprofiler-compute/docs/conceptual/cdna/l2-cache.rst)
+- [Local Data Share](https://github.com/ROCm/rocm-systems/blob/8cce7054f17a77272e1ba73ef78f4f77cf056956/projects/rocprofiler-compute/docs/conceptual/cdna/local-data-share.rst)
+- [Infinity Fabric transactions](https://github.com/ROCm/rocm-systems/blob/8cce7054f17a77272e1ba73ef78f4f77cf056956/projects/rocprofiler-compute/docs/tutorial/includes/infinity-fabric-transactions.rst)
+- [Vector memory operation counting](https://github.com/ROCm/rocm-systems/blob/8cce7054f17a77272e1ba73ef78f4f77cf056956/projects/rocprofiler-compute/docs/tutorial/includes/vector-memory-operation-counting.rst)
+- [LDS examples](https://github.com/ROCm/rocm-systems/blob/8cce7054f17a77272e1ba73ef78f4f77cf056956/projects/rocprofiler-compute/docs/tutorial/includes/lds-examples.rst)
+
+Client APUs use a different hierarchy; see the
+[RDNA performance model](https://github.com/ROCm/rocm-systems/blob/8cce7054f17a77272e1ba73ef78f4f77cf056956/projects/rocprofiler-compute/docs/conceptual/rdna/rdna-performance-model.rst)
+for GL0, GL1, and GL2.
+
+Read hit rates and utilizations against the `Percent of Peak` and `Peak`
+columns the report prints. Do not apply a remembered cutoff such as "L2 hit
+rate below 50% is bad"; a streaming kernel with no reuse is supposed to miss.
+Judge a rate against what the algorithm should do, the neighbouring kernels,
+and a known-good baseline.
+
+What usually follows from what:
+
+| Reading | Usual meaning |
+|---|---|
+| High request count, low cache hits | poor locality or no reuse to exploit |
+| Many requests per instruction in TA/TD | uncoalesced access |
+| Sustained LDS bank conflicts | padding or access reordering needed |
+| High fabric read or write stalls | traffic exceeds what the fabric can absorb |
+| Near-peak HBM with low cache hits | reduce traffic or add reuse; more bandwidth is not available |
+
+## 4. Memory bandwidth analysis (experimental, gfx950)
+
+A guided breakdown that walks the bandwidth tree and names the limiting level.
+Which GPUs it supports, and the experimental flag, are in
+[memory bandwidth analysis](https://github.com/ROCm/rocm-systems/blob/8cce7054f17a77272e1ba73ef78f4f77cf056956/projects/rocprofiler-compute/docs/how-to/membw_analysis.rst). It is
+experimental and currently shipped for gfx950 only. It needs only the
+memory chart (block 3) and block 30, so profile only those two blocks:
+
+```bash
+rocprof-compute profile --experimental --membw-analysis -b 3 30 \
+    --name <name> -- <application>
+
+rocprof-compute analyze --path ./workloads/<name>/<gpu_model> -k <kernel_id>
+```
+
+Its cutoffs and guidance text are data, not something to restate from memory.
+In a source checkout, read:
+
+- [gfx950_membw_tree_spec.yaml](https://github.com/ROCm/rocm-systems/blob/8cce7054f17a77272e1ba73ef78f4f77cf056956/projects/rocprofiler-compute/src/membw_analysis/tree_spec/gfx950_membw_tree_spec.yaml)
+- [gfx950_membw_guidance.yaml](https://github.com/ROCm/rocm-systems/blob/8cce7054f17a77272e1ba73ef78f4f77cf056956/projects/rocprofiler-compute/src/membw_analysis/tree_spec/gfx950_membw_guidance.yaml)
+
+A ROCm install copies the skills to `<rocm-path>/share/rocprofiler-compute/skills/`
+and does not copy these YAML files next to them. When this skill was installed
+rather than opened from the source tree, read those same two paths in the
+matching rocprofiler-compute checkout. Report the guidance the tool produced
+rather than substituting your own numbers.
+
+On any other architecture, use the Memory Chart path above.
+
+## 5. Normalization
+
+For a bandwidth question, `-n per_second` is meaningful and `per_cycle` is not.
+`-n` applies to the entire report, so switch it only when the user is asking
+about bandwidth alone. Leave it at `per_kernel` for a mixed report.
diff --git a/skills/rocprof-compute-memory/evals/evals.json b/skills/rocprof-compute-memory/evals/evals.json
new file mode 100644
index 00000000..0fb25faf
--- /dev/null
+++ b/skills/rocprof-compute-memory/evals/evals.json
@@ -0,0 +1,84 @@
+{
+  "evaluations": [
+    {
+      "id": "mem-access-inefficiency",
+      "skill_should_trigger": true,
+      "prompt": "find the memory access inefficiencies in this rocprofiler-compute capture\n\nThe application is ./vrandom_access.cpp.",
+      "workspace": "../rocprof-compute-sample",
+      "expected_behavior": [
+        "reads the memory chart before opening an individual cache level",
+        "cites hit rates or request counts from the report",
+        "ties any coalescing or tiling advice to those numbers"
+      ],
+      "unexpected_behavior": [
+        "recommends an optimization without citing a metric from the report"
+      ],
+      "logs_contain": ["-b memchart"]
+    },
+    {
+      "id": "mem-cache-hit-rates",
+      "skill_should_trigger": true,
+      "prompt": "what are my L2 and vL1D hit rates for this kernel and are they bad?\n\nThe application is ./vsequential_access.cpp.",
+      "workspace": "../rocprof-compute-sample",
+      "expected_behavior": [
+        "opens the l2 and vl1d blocks for a single kernel id",
+        "judges the rates against what the algorithm should do rather than a fixed cutoff"
+      ],
+      "unexpected_behavior": [
+        "applies a remembered threshold such as L2 below 50% being bad, as if the tool defined one"
+      ],
+      "logs_contain": ["-b"]
+    },
+    {
+      "id": "mem-lds-bank-conflicts",
+      "skill_should_trigger": true,
+      "prompt": "does my kernel have LDS bank conflicts?\n\nThe application is ./lds.hip.",
+      "workspace": "../rocprof-compute-sample",
+      "expected_behavior": [
+        "opens the lds block for the selected kernel",
+        "reads the bank conflict counters from the report"
+      ],
+      "logs_contain": ["lds"]
+    },
+    {
+      "id": "mem-bandwidth-bound",
+      "skill_should_trigger": true,
+      "prompt": "this kernel looks memory bound, where is the bandwidth going?\n\nThe application is ./vcopy.cpp.",
+      "workspace": "../rocprof-compute-sample",
+      "expected_behavior": [
+        "walks the hierarchy from the memory chart down to the level responsible",
+        "distinguishes traffic that caches absorbed from traffic that reached HBM"
+      ],
+      "logs_contain": ["memchart"]
+    },
+    {
+      "id": "mem-membw-analysis-gfx950",
+      "skill_should_trigger": true,
+      "prompt": "run the guided memory bandwidth analysis on my MI350 workload\n\nThe application is ./gl2_backpressure.hip.",
+      "workspace": "../rocprof-compute-sample/membw_analysis_test_suite",
+      "expected_behavior": [
+        "profiles only blocks 3 and 30 with --experimental --membw-analysis",
+        "notes that the analysis is experimental and shipped for gfx950",
+        "reports the guidance the tool produced rather than substituting its own cutoffs"
+      ],
+      "logs_contain": ["--membw-analysis"]
+    },
+    {
+      "id": "mem-host-memory",
+      "skill_should_trigger": false,
+      "note": "Memory question, but host-side allocation is not GPU counter analysis.",
+      "prompt": "my application is leaking host memory, help me find it"
+    },
+    {
+      "id": "mem-cuda-memory",
+      "skill_should_trigger": false,
+      "prompt": "analyze the memory throughput of my CUDA kernel with Nsight"
+    },
+    {
+      "id": "mem-gpu-oom",
+      "skill_should_trigger": false,
+      "note": "About allocation failure, not bandwidth or cache behaviour.",
+      "prompt": "why does my model run out of GPU memory during training?"
+    }
+  ]
+}
diff --git a/skills/rocprof-compute-memory/evals/machine.yml b/skills/rocprof-compute-memory/evals/machine.yml
new file mode 100644
index 00000000..4d7efd45
--- /dev/null
+++ b/skills/rocprof-compute-memory/evals/machine.yml
@@ -0,0 +1,2 @@
+os: [Linux]
+labels: [mi300x, gpu, rocm]
diff --git a/skills/rocprof-compute-memory/skill-card.md b/skills/rocprof-compute-memory/skill-card.md
new file mode 100644
index 00000000..6f0244bc
--- /dev/null
+++ b/skills/rocprof-compute-memory/skill-card.md
@@ -0,0 +1,13 @@
+# Skill Card
+
+## Description
+
+Analyzes the AMD GPU memory hierarchy with rocprof-compute, from the Memory Chart down to caches, LDS, and memory bandwidth.
+
+## Owner
+
+rocprofiler-compute, a project in [rocm-systems](https://github.com/ROCm/rocm-systems/)
+
+## License
+
+MIT
diff --git a/skills/rocprof-compute-pc-sampling/.federated.json b/skills/rocprof-compute-pc-sampling/.federated.json
new file mode 100644
index 00000000..a091cc85
--- /dev/null
+++ b/skills/rocprof-compute-pc-sampling/.federated.json
@@ -0,0 +1,10 @@
+{
+  "source": "rocm-rocm-systems",
+  "repo": "ROCm/rocm-systems",
+  "ref": "develop",
+  "commit": "8cce7054f17a77272e1ba73ef78f4f77cf056956",
+  "path": "projects/rocprofiler-compute/skills/rocprof-compute-pc-sampling",
+  "license": "MIT",
+  "content_hash": "0fccdde867de5668a0c767e64f6274714f07befea1e4f9a2299c6c580a0072c6",
+  "imported_at": "2026-10-28T01:19:39Z"
+}
diff --git a/skills/rocprof-compute-pc-sampling/SKILL.md b/skills/rocprof-compute-pc-sampling/SKILL.md
new file mode 100644
index 00000000..e3857ec8
--- /dev/null
+++ b/skills/rocprof-compute-pc-sampling/SKILL.md
@@ -0,0 +1,123 @@
+---
+name: rocprof-compute-pc-sampling
+description: Collects and analyzes rocprof-compute PC sampling to show which individual instructions in an AMD GPU kernel are hot and why wavefronts stall there, with ISA and source line attribution. Use when the user asks which instruction or source line is slow, asks about stall reasons, wants instruction-level or ISA-level detail, or wants hotspots inside a single kernel. Not for CUDA tools, Windows, or whole-application timing.
+---
+
+# PC sampling
+
+PC sampling answers a question counters cannot: which instruction inside the
+kernel is hot, and what the wavefront was waiting on. Use it after counters
+have already identified the kernel worth this level of detail.
+
+PC sampling is experimental. Every command below needs `--experimental`, and
+option names can change between releases, so run
+`rocprof-compute profile --experimental --help` before choosing flags.
+
+This skill drives `rocprof-compute --pc-sampling`. A request to invoke
+`rocprofv3 --pc-sampling-beta-enabled` directly belongs to the rocprofiler-sdk
+PC sampling skill, not this one. Method support is in
+[PC sampling](https://github.com/ROCm/rocm-systems/blob/8cce7054f17a77272e1ba73ef78f4f77cf056956/projects/rocprofiler-compute/docs/how-to/pc_sampling.rst).
+
+## 1. Check what the hardware supports
+
+```bash
+rocprofv3-avail info --pc-sampling
+```
+
+This reports the sampling methods the installed device accepts and the valid
+interval range. Two methods exist:
+
+- `stochastic` samples in hardware and is the only method that records a stall
+  reason. Which GPUs accept it is in
+  [PC sampling](https://github.com/ROCm/rocm-systems/blob/8cce7054f17a77272e1ba73ef78f4f77cf056956/projects/rocprofiler-compute/docs/how-to/pc_sampling.rst).
+- `host_trap` records no stall reason, and its samples can land a little past
+  the instruction responsible.
+
+`--pc-sampling-method` defaults to `stochastic`. Fall back to `host_trap` only
+when the device rejects stochastic.
+
+Never hardcode an interval. Stochastic intervals are counted in cycles and must
+be a power of two; host-trap intervals are microseconds. The accepted range
+comes from the device, so take it from the command above.
+
+## 2. Build for source attribution
+
+Samples always map to the kernel's assembly. Mapping assembly back to a source
+line requires debug info in the application, for example `hipcc -g`. Without
+it, `source_line` reads `N/A` and the `Source` column in CSV output is empty.
+Keep optimization flags on; profiling a `-O0` build measures the wrong program.
+
+Ask about the build before collecting, because rebuilding afterwards means
+re-profiling.
+
+## 3. Collect
+
+```bash
+rocprof-compute profile \
+    --experimental \
+    --pc-sampling \
+    --pc-sampling-method stochastic \
+    --pc-sampling-interval <interval-from-the-query> \
+    --name <workload_name> \
+    -- <application> <args>
+```
+
+PC sampling does not collect counters. A PC-sampling-only profile cannot be
+combined with the framework tracing options, and it will not populate the
+counter panels. When the user needs both, collect two workloads.
+
+## 4. Analyze
+
+```bash
+rocprof-compute analyze --path ./workloads/<name>/<gpu_model> --list-stats
+
+rocprof-compute analyze \
+    --path ./workloads/<name>/<gpu_model> \
+    -k <kernel_id> \
+    --pc-sampling-sorting-type count \
+    --pc-sampling-rows 20
+```
+
+Sorting defaults to `count`, which puts the hottest instruction first. Use
+`offset` to read the kernel in program order instead. `--pc-sampling-rows`
+defaults to 10; pass `0` for every row.
+
+Read the table as:
+
+| Column | Meaning |
+|---|---|
+| PC offset | where the instruction sits in the kernel's ISA |
+| Sample count | how often execution was there |
+| Stall reason | what the wavefront was waiting on, stochastic sampling only |
+
+A high sample count marks a hot instruction, not a defective one. A loop body
+is supposed to be hot. What matters is a hot instruction whose stall reason
+says it is waiting: a memory stall on a load means the kernel is latency-bound
+there, and the fix is usually coalescing, prefetching, or more independent work
+between the load and its use.
+
+Stall reasons and what each one implies are listed in
+[pc_sampling.rst](https://github.com/ROCm/rocm-systems/blob/8cce7054f17a77272e1ba73ef78f4f77cf056956/projects/rocprofiler-compute/docs/how-to/pc_sampling.rst). Do not interpret a stall
+reason from host-trap data; it does not record one.
+
+## 5. Export annotated disassembly
+
+```bash
+rocprof-compute analyze --path ./workloads/<name>/<gpu_model> \
+    --output-format csv --output-name pc_report
+```
+
+CSV output writes per-kernel annotated disassembly into a
+`per_kernel_pc_sampling/` directory, with sample counts and source lines beside
+each instruction. `--output-format db` stores the same content in the analysis
+database. Both disable terminal output.
+
+## 6. When there are too few samples
+
+A short kernel may not accumulate enough samples to be meaningful. Increase the
+work the kernel does, or shorten the interval within the range the device
+accepts, rather than drawing conclusions from a handful of samples.
+
+For multi-process workloads, keep each process or rank separate; see the
+multi-process sections of
+[pc_sampling.rst](https://github.com/ROCm/rocm-systems/blob/8cce7054f17a77272e1ba73ef78f4f77cf056956/projects/rocprofiler-compute/docs/how-to/pc_sampling.rst).
diff --git a/skills/rocprof-compute-pc-sampling/evals/evals.json b/skills/rocprof-compute-pc-sampling/evals/evals.json
new file mode 100644
index 00000000..d9aa23a8
--- /dev/null
+++ b/skills/rocprof-compute-pc-sampling/evals/evals.json
@@ -0,0 +1,82 @@
+{
+  "evaluations": [
+    {
+      "id": "pcs-which-instruction-stalls",
+      "skill_should_trigger": true,
+      "prompt": "which instruction in my MI300 kernel stalls the most?\n\nThe application is ./vmem.hip.",
+      "workspace": "../rocprof-compute-sample",
+      "expected_behavior": [
+        "queries the device for supported PC sampling configurations before choosing an interval",
+        "collects with --experimental --pc-sampling",
+        "sorts the analysis by sample count and reads the stall reason column"
+      ],
+      "unexpected_behavior": [
+        "hardcodes a sampling interval without querying the device"
+      ],
+      "logs_contain": ["--pc-sampling", "rocprofv3-avail"]
+    },
+    {
+      "id": "pcs-source-line-attribution",
+      "skill_should_trigger": true,
+      "prompt": "map the hot GPU instructions back to my HIP source lines\n\nThe application is ./vmem.hip.",
+      "workspace": "../rocprof-compute-sample",
+      "expected_behavior": [
+        "states that source line mapping requires the application to be built with debug info",
+        "notes that samples map to assembly even without debug info"
+      ],
+      "unexpected_behavior": [
+        "promises source line attribution for a build with no debug info"
+      ]
+    },
+    {
+      "id": "pcs-stochastic-vs-host-trap",
+      "skill_should_trigger": true,
+      "prompt": "should I use stochastic or host trap PC sampling on this GPU?",
+      "expected_behavior": [
+        "explains that only stochastic records a stall reason",
+        "notes stochastic needs MI300 or newer and host_trap is the fallback",
+        "checks what the installed device actually supports"
+      ],
+      "logs_contain": ["stochastic"]
+    },
+    {
+      "id": "pcs-export-disassembly",
+      "skill_should_trigger": true,
+      "prompt": "export the annotated disassembly with sample counts from my pc sampling workload\n\nThe application is ./vmem.hip.",
+      "workspace": "../rocprof-compute-sample",
+      "expected_behavior": [
+        "uses --output-format csv or db",
+        "notes that the per-kernel disassembly is written under per_kernel_pc_sampling"
+      ],
+      "logs_contain": ["--output-format"]
+    },
+    {
+      "id": "pcs-host-trap-stall-reason",
+      "skill_should_trigger": true,
+      "prompt": "my host trap PC sampling run has no stall reasons, how do I interpret them?",
+      "expected_behavior": [
+        "explains that host_trap records no stall reason",
+        "offers stochastic sampling if the device supports it"
+      ],
+      "unexpected_behavior": [
+        "interprets a stall reason from host-trap data"
+      ]
+    },
+    {
+      "id": "pcs-rank-kernels-by-time",
+      "skill_should_trigger": false,
+      "note": "Hotspot wording, but ranking whole kernels is counter work, not instruction sampling.",
+      "prompt": "which of my GPU kernels accounts for the most total time?"
+    },
+    {
+      "id": "pcs-cpu-profiler",
+      "skill_should_trigger": false,
+      "prompt": "profile my C++ program's CPU hotspots with perf"
+    },
+    {
+      "id": "pcs-cuda-sampling",
+      "skill_should_trigger": false,
+      "prompt": "do instruction level sampling on my CUDA kernel"
+    }
+  ]
+}
diff --git a/skills/rocprof-compute-pc-sampling/evals/machine.yml b/skills/rocprof-compute-pc-sampling/evals/machine.yml
new file mode 100644
index 00000000..4d7efd45
--- /dev/null
+++ b/skills/rocprof-compute-pc-sampling/evals/machine.yml
@@ -0,0 +1,2 @@
+os: [Linux]
+labels: [mi300x, gpu, rocm]
diff --git a/skills/rocprof-compute-pc-sampling/skill-card.md b/skills/rocprof-compute-pc-sampling/skill-card.md
new file mode 100644
index 00000000..bf8bb891
--- /dev/null
+++ b/skills/rocprof-compute-pc-sampling/skill-card.md
@@ -0,0 +1,13 @@
+# Skill Card
+
+## Description
+
+Uses rocprof-compute PC sampling to show which instructions in an AMD GPU kernel are hot and why they stall.
+
+## Owner
+
+rocprofiler-compute, a project in [rocm-systems](https://github.com/ROCm/rocm-systems/)
+
+## License
+
+MIT
diff --git a/skills/rocprof-compute-roofline/.federated.json b/skills/rocprof-compute-roofline/.federated.json
new file mode 100644
index 00000000..29ef0168
--- /dev/null
+++ b/skills/rocprof-compute-roofline/.federated.json
@@ -0,0 +1,10 @@
+{
+  "source": "rocm-rocm-systems",
+  "repo": "ROCm/rocm-systems",
+  "ref": "develop",
+  "commit": "8cce7054f17a77272e1ba73ef78f4f77cf056956",
+  "path": "projects/rocprofiler-compute/skills/rocprof-compute-roofline",
+  "license": "MIT",
+  "content_hash": "6624d509ddee66144bf15f8bb39d0ccfde81a0e536cf11e34dfed618d05c17eb",
+  "imported_at": "2026-10-28T01:19:39Z"
+}
diff --git a/skills/rocprof-compute-roofline/SKILL.md b/skills/rocprof-compute-roofline/SKILL.md
new file mode 100644
index 00000000..14eff91e
--- /dev/null
+++ b/skills/rocprof-compute-roofline/SKILL.md
@@ -0,0 +1,108 @@
+---
+name: rocprof-compute-roofline
+description: Collects and reads rocprof-compute roofline data to place an AMD GPU kernel against the hardware's peak compute and peak bandwidth ceilings. Use when the user asks whether a kernel is compute-bound or memory-bound, asks about arithmetic intensity or FLOPs per byte, asks how much performance is left on the table, or asks for a roofline. Not for CUDA tools, Windows, or per-instruction analysis.
+---
+
+# Roofline analysis
+
+A roofline places a kernel's arithmetic intensity against two measured
+ceilings, peak compute and peak bandwidth, and answers one question: which of
+the two is the limit.
+
+Run `rocprof-compute profile --help` and `rocprof-compute analyze --help`
+before choosing flags. Do not open the generated HTML chart. The analyze table
+carries the same numbers in readable form.
+
+Roofline support by GPU is in
+[supported hardware](https://github.com/ROCm/rocm-systems/blob/8cce7054f17a77272e1ba73ef78f4f77cf056956/projects/rocprofiler-compute/docs/what-is-rocprof-compute.rst).
+
+## 1. Collect
+
+Roofline is collected by default, so an ordinary profile already has it:
+
+```bash
+rocprof-compute profile --name <workload_name> -- <application> <args>
+```
+
+Collecting it runs microbenchmarks in addition to profiling the application,
+which takes extra time. `--no-roof` skips them when the user only wants
+counters.
+
+On MI100 (gfx908), the roofline microbenchmarks are skipped and the profile
+still succeeds. When [supported hardware](https://github.com/ROCm/rocm-systems/blob/8cce7054f17a77272e1ba73ef78f4f77cf056956/projects/rocprofiler-compute/docs/what-is-rocprof-compute.rst)
+shows no roofline support, do not ask the user to re-profile without
+`--no-roof`. Use Speed-of-Light and the `rocprof-compute-memory` skill instead.
+
+`--device` and `HIP_VISIBLE_DEVICES` are different selections. See
+[Selecting a GPU](https://github.com/ROCm/rocm-systems/blob/8cce7054f17a77272e1ba73ef78f4f77cf056956/projects/rocprofiler-compute/docs/how-to/profile/mode.rst).
+
+See [standalone roofline](https://github.com/ROCm/rocm-systems/blob/8cce7054f17a77272e1ba73ef78f4f77cf056956/projects/rocprofiler-compute/docs/how-to/profile/mode.rst) for `--roof-only`
+and `--bench-only`, which run the microbenchmarks without profiling an
+application.
+
+## 2. Read the table
+
+```bash
+rocprof-compute analyze --path ./workloads/<name>/<gpu_model> --list-stats
+rocprof-compute analyze --path ./workloads/<name>/<gpu_model> -k <kernel_id> -b roof
+```
+
+Always pass one kernel id. The table gives the kernel's arithmetic intensity in
+FLOPs per byte alongside the empirical peak FLOPs and peak bandwidth measured
+on this machine. Compare against those printed numbers; never compute a peak
+yourself.
+
+Report two things first, before any tuning advice:
+
+1. **Which ceiling limits the kernel**, memory or compute.
+2. **How far the kernel sits below that ceiling**, as a percentage of it.
+
+Read the verdict from the table as:
+
+- Low arithmetic intensity, close to the bandwidth ceiling: memory-bound. Work
+  on traffic, reuse, tiling, and locality. Continue with the
+  `rocprof-compute-memory` skill.
+- High arithmetic intensity, close to the compute ceiling: compute-bound. Work
+  on instruction mix, matrix instruction use, and divergence.
+- Well below both ceilings: neither is the limit. The kernel is held back by
+  occupancy, launch configuration, or stalls. Go back to the
+  `rocprof-compute-kernel-bottleneck` skill.
+
+A kernel sitting below both ceilings is the common case and the most
+misread one. Never report it as compute-bound merely because its arithmetic
+intensity is high.
+
+The distance to the ceiling is what makes the verdict actionable. "Memory-bound
+at 45% of peak bandwidth" tells the user there is headroom; "memory-bound at
+94%" tells them to reduce traffic instead of chasing it.
+
+## 3. Pick the right precision and memory level
+
+The table lists a peak for every precision and an arithmetic intensity for
+every memory level. Compare against the peak row for the precision the kernel
+actually issues, and the intensity row for the memory level in question.
+
+A mixed-precision or matrix kernel compared against the FP32 peak will look far
+worse than it is. Read the precision from the instruction mix with `-b cu_ins`
+rather than assuming it. Ask the user which precision the kernel is meant to
+use only when the mix is genuinely mixed and no single type dominates; they
+know the intent, the counters know what ran.
+
+HBM answers "am I limited by main memory", while L2 and vL1D show whether
+caches are absorbing the traffic.
+
+`--view table` prints one table for all kernels together instead of one per
+kernel. To compare many kernels, write the
+[analysis database](https://github.com/ROCm/rocm-systems/blob/8cce7054f17a77272e1ba73ef78f4f77cf056956/projects/rocprofiler-compute/docs/how-to/analyze/cli.rst) with
+`--output-format db` and query it with SQL.
+
+## 4. When roofline data is missing
+
+| Cause | What to do |
+|---|---|
+| Profiled with `--no-roof` | re-profile without it |
+| Architecture has no microbenchmark support | use Speed-of-Light and the `rocprof-compute-memory` skill instead |
+| `roofline.csv` absent from the workload directory | the benchmark did not complete; check the profile log |
+
+The concepts behind the ceilings are in the
+[performance model](https://github.com/ROCm/rocm-systems/blob/8cce7054f17a77272e1ba73ef78f4f77cf056956/projects/rocprofiler-compute/docs/conceptual/performance-model.rst).
diff --git a/skills/rocprof-compute-roofline/evals/evals.json b/skills/rocprof-compute-roofline/evals/evals.json
new file mode 100644
index 00000000..23a8e7bd
--- /dev/null
+++ b/skills/rocprof-compute-roofline/evals/evals.json
@@ -0,0 +1,73 @@
+{
+  "evaluations": [
+    {
+      "id": "roof-compute-or-memory-bound",
+      "skill_should_trigger": true,
+      "prompt": "is this rocprof-compute kernel compute bound or memory bound?\n\nThe application is ./vcopy.cpp.",
+      "workspace": "../rocprof-compute-sample",
+      "expected_behavior": [
+        "analyzes the roof block for a single kernel id",
+        "compares arithmetic intensity against the empirical peaks printed in the table",
+        "says when the kernel is below both ceilings instead of forcing a bound"
+      ],
+      "unexpected_behavior": [
+        "opens the generated roofline HTML chart instead of reading the table",
+        "calculates a peak instead of using the measured one"
+      ],
+      "logs_contain": ["-b roof"]
+    },
+    {
+      "id": "roof-arithmetic-intensity",
+      "skill_should_trigger": true,
+      "prompt": "what is the arithmetic intensity of my hottest kernel and how much performance am I leaving on the table?\n\nThe application is ./mat_mul_max.hip.",
+      "workspace": "../rocprof-compute-sample",
+      "expected_behavior": [
+        "reads FLOPs per byte and the peak values from the roofline table",
+        "narrows to one kernel id first"
+      ],
+      "logs_contain": ["roof"]
+    },
+    {
+      "id": "roof-precision-selection",
+      "skill_should_trigger": true,
+      "prompt": "my half precision GEMM looks terrible on the roofline, is that right?\n\nThe application is ./wmma_gemm_fp16.hip.",
+      "workspace": "../rocprof-compute-sample/wmma_gemm_fp16",
+      "expected_behavior": [
+        "compares against the peak for the precision the kernel actually issues instead of FP32",
+        "confirms the precision from the instruction mix rather than assuming it"
+      ],
+      "unexpected_behavior": [
+        "compares a half precision kernel against the FP32 ceiling and calls it slow"
+      ],
+      "logs_contain": ["cu_ins"]
+    },
+    {
+      "id": "roof-missing-data",
+      "skill_should_trigger": true,
+      "prompt": "my rocprof-compute report has no roofline section, what went wrong?",
+      "expected_behavior": [
+        "checks whether the profile used --no-roof",
+        "checks whether the architecture supports the roofline microbenchmarks before advising a re-profile"
+      ],
+      "unexpected_behavior": [
+        "tells the user to re-profile without --no-roof when the architecture is the actual reason"
+      ]
+    },
+    {
+      "id": "roof-generic-optimization",
+      "skill_should_trigger": false,
+      "prompt": "how can I make my application faster?"
+    },
+    {
+      "id": "roof-cuda-roofline",
+      "skill_should_trigger": false,
+      "prompt": "generate a roofline for my CUDA kernel with Nsight Compute"
+    },
+    {
+      "id": "roof-concept-question",
+      "skill_should_trigger": false,
+      "note": "Asks what a roofline is in the abstract, with no workload to analyze.",
+      "prompt": "explain the roofline performance model to me"
+    }
+  ]
+}
diff --git a/skills/rocprof-compute-roofline/evals/machine.yml b/skills/rocprof-compute-roofline/evals/machine.yml
new file mode 100644
index 00000000..4d7efd45
--- /dev/null
+++ b/skills/rocprof-compute-roofline/evals/machine.yml
@@ -0,0 +1,2 @@
+os: [Linux]
+labels: [mi300x, gpu, rocm]
diff --git a/skills/rocprof-compute-roofline/skill-card.md b/skills/rocprof-compute-roofline/skill-card.md
new file mode 100644
index 00000000..120ea14f
--- /dev/null
+++ b/skills/rocprof-compute-roofline/skill-card.md
@@ -0,0 +1,13 @@
+# Skill Card
+
+## Description
+
+Places an AMD GPU kernel against measured peak compute and peak bandwidth ceilings with rocprof-compute roofline data.
+
+## Owner
+
+rocprofiler-compute, a project in [rocm-systems](https://github.com/ROCm/rocm-systems/)
+
+## License
+
+MIT
diff --git a/skills/rocprof-compute-sample/.fixture.json b/skills/rocprof-compute-sample/.fixture.json
new file mode 100644
index 00000000..670da40c
--- /dev/null
+++ b/skills/rocprof-compute-sample/.fixture.json
@@ -0,0 +1,8 @@
+{
+  "repo": "ROCm/rocm-systems",
+  "ref": "develop",
+  "commit": "8cce7054f17a77272e1ba73ef78f4f77cf056956",
+  "path": "projects/rocprofiler-compute/sample",
+  "content_hash": "20210259dfbeada334ea1a051d146f05fde13a280ef6ac4523931f725f66f001",
+  "imported_at": "2026-10-28T01:19:39Z"
+}
diff --git a/skills/rocprof-compute-sample/common.h b/skills/rocprof-compute-sample/common.h
new file mode 100644
index 00000000..2609685a
--- /dev/null
+++ b/skills/rocprof-compute-sample/common.h
@@ -0,0 +1,20 @@
+// Copyright (c) Advanced Micro Devices, Inc.
+// SPDX-License-Identifier:  MIT
+
+#pragma once
+
+#include <hip/hip_runtime.h>
+#include <iostream>
+
+#define hipCheck(stmt)                                                         \
+  do {                                                                         \
+    hipError_t err = stmt;                                                     \
+    if (err != hipSuccess) {                                                   \
+      char msg[256];                                                           \
+      sprintf(msg, "%s in file %s, function %s, line %d\n", #stmt, __FILE__,   \
+              __FUNCTION__, __LINE__);                                         \
+      std::string errstring = hipGetErrorString(err);                          \
+      std::cerr << msg << "\t" << errstring << std::endl;                      \
+      throw std::runtime_error(msg);                                           \
+    }                                                                          \
+  } while (0)
diff --git a/skills/rocprof-compute-sample/lds.hip b/skills/rocprof-compute-sample/lds.hip
new file mode 100644
index 00000000..887d4b5d
--- /dev/null
+++ b/skills/rocprof-compute-sample/lds.hip
@@ -0,0 +1,54 @@
+// Copyright (c) Advanced Micro Devices, Inc.
+// SPDX-License-Identifier:  MIT
+
+/*
+An example to explore LDS bandwidth and bank conflicts, written by Nicholas Curtis [AMD].
+*/
+
+#include "common.h"
+
+constexpr unsigned max_threads = 256;
+constexpr unsigned nbanks = 32;
+
+__global__ void load(int* out, int flag) {
+  __shared__ int array[max_threads];
+  int index = threadIdx.x;
+  // fake a store to the LDS array to avoid unwanted behavior
+  if (flag)
+    array[max_threads - index] = index;
+  __syncthreads();
+  int x = array[index];
+  if (x == int(-1234567))
+    out[threadIdx.x] = x;
+}
+
+__global__ void conflicts(int* out, int flag) {
+  constexpr unsigned nelements = nbanks * max_threads;
+  __shared__ int array[nelements];
+  // each thread reads from the same bank
+  int index = threadIdx.x * nbanks;
+  // fake a store to the LDS array to avoid unwanted behavior
+  if (flag)
+    array[max_threads - index] = index;
+  __syncthreads();
+  int x = array[index];
+  if (x == int(-1234567))
+    out[threadIdx.x] = x;
+}
+
+void bandwidth_demo(int N) {
+  for (int i = 1; i <= N; ++i)
+    load<<<1,i>>>(nullptr, 0);
+  hipCheck(hipDeviceSynchronize());
+}
+
+void conflicts_demo(int N) {
+  for (int i = 1; i <= N; ++i)
+    conflicts<<<1,i>>>(nullptr, 0);
+  hipCheck(hipDeviceSynchronize());
+}
+
+int main() {
+  bandwidth_demo(max_threads);
+  conflicts_demo(max_threads);
+}
diff --git a/skills/rocprof-compute-sample/mat_mul_max.hip b/skills/rocprof-compute-sample/mat_mul_max.hip
new file mode 100644
index 00000000..2a9791c4
--- /dev/null
+++ b/skills/rocprof-compute-sample/mat_mul_max.hip
@@ -0,0 +1,92 @@
+// Copyright (c) Advanced Micro Devices, Inc.
+// SPDX-License-Identifier:  MIT
+
+#include <hip/hip_runtime.h>
+#include <iostream>
+#include <cstdlib>
+#include <chrono>
+
+#define TILE_SIZE 32     // Maximum block size: 32 x 32 = 1024 threads/block
+#define N 4096           // Matrix size: 4096 x 4096 (~67M elements)
+
+// Helper macro for HIP error checking
+#define HIP_CHECK(call)                                              \
+    do {                                                            \
+        hipError_t err = call;                                      \
+        if (err != hipSuccess) {                                    \
+            std::cerr << "HIP error: " << hipGetErrorString(err)    \
+                      << " at " << __FILE__ << ":" << __LINE__      \
+                      << std::endl;                                 \
+            std::exit(EXIT_FAILURE);                                \
+        }                                                           \
+    } while(0)
+
+__global__ void matMulKernel(const float* __restrict__ A, const float* __restrict__ B, float* __restrict__ C, int width) {
+    __shared__ float tileA[TILE_SIZE][TILE_SIZE];
+    __shared__ float tileB[TILE_SIZE][TILE_SIZE];
+
+    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
+    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
+
+    float sum = 0.0f;
+
+    // Loop over tiles of input matrices
+    for (int t = 0; t < width / TILE_SIZE; ++t) {
+        tileA[threadIdx.y][threadIdx.x] = A[row * width + t * TILE_SIZE + threadIdx.x];
+        tileB[threadIdx.y][threadIdx.x] = B[(t * TILE_SIZE + threadIdx.y) * width + col];
+        __syncthreads();
+
+        #pragma unroll
+        for (int i = 0; i < TILE_SIZE; ++i)
+            sum += tileA[threadIdx.y][i] * tileB[i][threadIdx.x];
+
+        __syncthreads();
+    }
+
+    C[row * width + col] = sum;
+}
+
+int main() {
+    size_t size = N * N * sizeof(float);
+
+    float *h_A = new float[N * N];
+    float *h_B = new float[N * N];
+    float *d_A, *d_B, *d_C;
+
+    // Initialize matrices with dummy values
+    for (int i = 0; i < N * N; ++i) {
+        h_A[i] = static_cast<float>(i % 100) * 0.01f;
+        h_B[i] = static_cast<float>((i + 1) % 100) * 0.01f;
+    }
+
+    HIP_CHECK(hipMalloc(&d_A, size));
+    HIP_CHECK(hipMalloc(&d_B, size));
+    HIP_CHECK(hipMalloc(&d_C, size));
+
+    HIP_CHECK(hipMemcpy(d_A, h_A, size, hipMemcpyHostToDevice));
+    HIP_CHECK(hipMemcpy(d_B, h_B, size, hipMemcpyHostToDevice));
+
+    dim3 blockDim(TILE_SIZE, TILE_SIZE);  // 32 x 32 = 1024 threads
+    dim3 gridDim(N / TILE_SIZE, N / TILE_SIZE);  // 128 x 128 = 16,384 thread blocks
+
+    std::cout << "Launching kernel with grid: (" << gridDim.x << ", " << gridDim.y
+              << ") and block: (" << blockDim.x << ", " << blockDim.y << ")\n";
+
+    auto start = std::chrono::high_resolution_clock::now();
+
+    matMulKernel<<<gridDim, blockDim>>>(d_A, d_B, d_C, N);
+    HIP_CHECK(hipGetLastError());
+    HIP_CHECK(hipDeviceSynchronize());
+
+    auto end = std::chrono::high_resolution_clock::now();
+    std::chrono::duration<double> elapsed = end - start;
+    std::cout << "Execution time: " << elapsed.count() << " seconds\n";
+
+    HIP_CHECK(hipFree(d_A));
+    HIP_CHECK(hipFree(d_B));
+    HIP_CHECK(hipFree(d_C));
+    delete[] h_A;
+    delete[] h_B;
+
+    return 0;
+}
diff --git a/skills/rocprof-compute-sample/membw_analysis_test_suite/gl2_backpressure.hip b/skills/rocprof-compute-sample/membw_analysis_test_suite/gl2_backpressure.hip
new file mode 100644
index 00000000..223b4295
--- /dev/null
+++ b/skills/rocprof-compute-sample/membw_analysis_test_suite/gl2_backpressure.hip
@@ -0,0 +1,96 @@
+// Copyright (c) Advanced Micro Devices, Inc.
+// SPDX-License-Identifier:  MIT
+
+#include <hip/hip_runtime.h>
+#include <iostream>
+#include <chrono>
+
+#define CHECK_HIP(call) \
+    do { \
+        hipError_t err = call; \
+        if (err != hipSuccess) { \
+            fprintf(stderr, "HIP error at %s:%d - %s\n", __FILE__, __LINE__, hipGetErrorString(err)); \
+            fflush(stderr); \
+            exit(1); \
+        } \
+    } while(0)
+
+constexpr long long L2_STRIDE = 67 * 1024 * 1024LL;  // 67MB - doesn't divide N
+
+// Baseline: Uncoalesced + large stride -> L2 misses + high request rate -> fill LFIFO -> TCR stall
+__global__ void gl2_backpressure_baseline(float* data, int N) {
+    int idx = blockIdx.x * blockDim.x + threadIdx.x;
+    if (idx >= N) return;
+
+    int lane = threadIdx.x & 63;  // Wave64
+
+    // UNCOALESCED
+    // LARGE STRIDE
+    float acc = 0.0f;
+    for (int i = 0; i < 128; i++) {
+        long long offset = (long long)idx * 64 + (long long)lane * 32 + (long long)i * L2_STRIDE;
+        int access_idx = (int)(offset % (long long)N);
+        acc += data[access_idx];
+    }
+    data[idx] = acc;
+}
+
+// Optimized: Shared memory, minimal L2 traffic, no write pressure
+__global__ void gl2_backpressure_optimized(float* data, int N) {
+    int idx = blockIdx.x * blockDim.x + threadIdx.x;
+    if (idx >= N) return;
+
+    __shared__ float sdata[256];
+    sdata[threadIdx.x] = data[idx];
+    __syncthreads();
+
+    float sum = sdata[threadIdx.x];
+    for (int i = 0; i < 64; i++) {
+        sum = sum * 0.99f + 1.0f;
+    }
+
+    data[idx] = sum;  // Single coalesced write
+}
+
+int main(int argc, char** argv) {
+    bool use_optimized = (argc > 1 && std::string(argv[1]) == "opt");
+
+    // 2GB
+    int N = 512 * 1024 * 1024;
+    size_t bytes = (size_t)N * sizeof(float);
+
+    float* d_data;
+    hipError_t alloc_err = hipMalloc(&d_data, bytes);
+    if (alloc_err != hipSuccess) {
+        // Fallback to 1GB if 2GB fails (still exceeds 64MB L2)
+        N = 256 * 1024 * 1024;
+        bytes = (size_t)N * sizeof(float);
+        CHECK_HIP(hipMalloc(&d_data, bytes));
+    }
+
+    // Host init + copy (no hipMemset - avoids __amd_rocclr_fillBufferAligned)
+    float* h_data = new float[N];
+    for (int i = 0; i < N; i++) h_data[i] = i * 0.001f;
+    CHECK_HIP(hipMemcpy(d_data, h_data, bytes, hipMemcpyHostToDevice));
+    delete[] h_data;
+
+    // Block 64 -> more blocks -> more concurrent waves
+    dim3 block(64);
+    dim3 grid((N + block.x - 1) / block.x);
+
+    // Single dispatch for clean profiling (no warmup to avoid double-dispatch in results)
+    auto start = std::chrono::high_resolution_clock::now();
+    if (use_optimized) {
+        hipLaunchKernelGGL(gl2_backpressure_optimized, grid, block, 0, 0, d_data, N);
+    } else {
+        hipLaunchKernelGGL(gl2_backpressure_baseline, grid, block, 0, 0, d_data, N);
+    }
+    CHECK_HIP(hipDeviceSynchronize());
+    auto end = std::chrono::high_resolution_clock::now();
+
+    std::cout << (use_optimized ? "Optimized" : "Baseline") << ": "
+              << std::chrono::duration_cast<std::chrono::milliseconds>(end - start).count() << " ms\n";
+
+    CHECK_HIP(hipFree(d_data));
+    return 0;
+}
diff --git a/skills/rocprof-compute-sample/occupancy.hip b/skills/rocprof-compute-sample/occupancy.hip
new file mode 100644
index 00000000..1c563203
--- /dev/null
+++ b/skills/rocprof-compute-sample/occupancy.hip
@@ -0,0 +1,85 @@
+// Copyright (c) Advanced Micro Devices, Inc.
+// SPDX-License-Identifier:  MIT
+
+/*
+An example to explore achieved occupancy, and various occupancy limiters.
+Written by Nicholas Curtis [AMD].
+*/
+
+#include "common.h"
+
+__global__ void empty(int N, double* ptr) {
+
+}
+
+constexpr int bound = 16;
+__launch_bounds__(256)
+__global__ void vgprbound(int N, double* ptr) {
+    double intermediates[bound];
+    for (int i = 0 ; i < bound; ++i) intermediates[i] = N * threadIdx.x;
+    double x = ptr[threadIdx.x];
+    for (int i = 0; i < 100; ++i) {
+        x += sin(pow(__shfl(x, i % warpSize) * intermediates[(i - 1) % bound], intermediates[i % bound]));
+        intermediates[i % bound] = x;
+    }
+    if (x == N) ptr[threadIdx.x] = x;
+}
+
+constexpr size_t fully_allocate_lds = 64ul * 1024ul / sizeof(double);
+__launch_bounds__(256)
+__global__ void ldsbound(int N, double* ptr) {
+    __shared__ double intermediates[fully_allocate_lds];
+    for (int i = threadIdx.x ; i < fully_allocate_lds; i += blockDim.x) intermediates[i] = N * threadIdx.x;
+    __syncthreads();
+    double x = ptr[threadIdx.x];
+    for (int i = threadIdx.x; i < fully_allocate_lds; i += blockDim.x) {
+        x += sin(pow(__shfl(x, i % warpSize) * intermediates[(i - 1) % fully_allocate_lds], intermediates[i % fully_allocate_lds]));
+        __syncthreads();
+        intermediates[i % fully_allocate_lds] = x;
+    }
+    if (x == N) ptr[threadIdx.x] = x;
+}
+
+constexpr int sgprlim = 1;
+__launch_bounds__(1024, 8)
+__global__ void sgprbound(int N, double* ptr) {
+    double intermediates[sgprlim];
+    for (int i = 0 ; i < sgprlim; ++i) intermediates[i] = i;
+    double x = ptr[0];
+    #pragma unroll 1
+    for (int i = 0; i < 100; ++i) {
+        x += sin(pow(intermediates[(i - 1) % sgprlim], intermediates[i % sgprlim]));
+        intermediates[i % sgprlim] = x;
+    }
+    if (x == N) ptr[0] = x;
+}
+
+int main() {
+    double* ptr;
+    hipCheck(hipMalloc(&ptr, 1024 * sizeof(double)));
+    vgprbound<<<1024 * 1024, 256>>>(0, ptr);
+    hipCheck(hipGetLastError());
+    hipCheck(hipDeviceSynchronize());
+    vgprbound<<<1024 * 1024, 256>>>(0, ptr);
+    hipCheck(hipGetLastError());
+    hipCheck(hipDeviceSynchronize());
+    ldsbound<<<1024 * 1024, 256>>>(0, ptr);
+    hipCheck(hipGetLastError());
+    hipCheck(hipDeviceSynchronize());
+    ldsbound<<<1024 * 1024, 256>>>(0, ptr);
+    hipCheck(hipGetLastError());
+    hipCheck(hipDeviceSynchronize());
+    sgprbound<<<1024 * 1024, 256>>>(0, ptr);
+    hipCheck(hipGetLastError());
+    hipCheck(hipDeviceSynchronize());
+    sgprbound<<<1024 * 1024, 256>>>(0, ptr);
+    hipCheck(hipGetLastError());
+    hipCheck(hipDeviceSynchronize());
+    empty<<<1024 * 1024, 256>>>(0, ptr);
+    hipCheck(hipGetLastError());
+    hipCheck(hipDeviceSynchronize());
+    empty<<<1024 * 1024, 256>>>(0, ptr);
+    hipCheck(hipGetLastError());
+    hipCheck(hipDeviceSynchronize());
+    hipCheck(hipFree(ptr));
+}
diff --git a/skills/rocprof-compute-sample/simple_net.py b/skills/rocprof-compute-sample/simple_net.py
new file mode 100755
index 00000000..6def8f28
--- /dev/null
+++ b/skills/rocprof-compute-sample/simple_net.py
@@ -0,0 +1,127 @@
+# Copyright (c) Advanced Micro Devices, Inc.
+# SPDX-License-Identifier:  MIT
+
+"""
+Sample PyTorch workload for rocprof-compute --torch-trace.
+
+    rocprof-compute profile --experimental --torch-trace --no-roof \
+        -n simple_net -- python3 ./simple_net.py
+
+    --user-range         Wrap the run in a user ROCTX range. Requires ROCTX Python.
+    --same-line-linear   Two Linear.forward calls from the same source line.
+    --no-backward        Skip loss.backward(). Cannot combine with --backward-thread.
+    --backward-thread    Run loss.backward() on a worker thread. Cannot combine
+                         with --no-backward.
+"""
+
+import argparse
+import sys
+import threading
+
+import torch
+import torch.nn as nn
+import torch.nn.functional as F
+
+try:
+    import roctx
+except ImportError:
+    roctx = None
+
+
+class SimpleNet(nn.Module):
+    def __init__(self, same_line_linear=False):
+        super().__init__()
+        self.fc1 = nn.Linear(10, 20)
+        self.fc2 = nn.Linear(20, 10)
+        self.same_line_linear = same_line_linear
+
+    def forward(self, x):
+        if self.same_line_linear:
+            # two Linear.forward calls on one source line
+            # fmt: off
+            x = self.fc1(x); x = self.fc2(x)  # noqa: E702
+            # fmt: on
+            return x
+        x = self.fc1(x)
+        x = F.relu(x)
+        x = self.fc2(x)
+        return x
+
+
+def parse_args(argv=None):
+    parser = argparse.ArgumentParser(description="SimpleNet torch-trace sample")
+    parser.add_argument(
+        "--user-range",
+        action="store_true",
+        help="Wrap the run in a user ROCTX range",
+    )
+    parser.add_argument(
+        "--same-line-linear",
+        action="store_true",
+        help="Two Linear.forward calls from the same source line",
+    )
+    exclusive = parser.add_mutually_exclusive_group()
+    exclusive.add_argument(
+        "--no-backward",
+        action="store_true",
+        help="Skip loss.backward()",
+    )
+    exclusive.add_argument(
+        "--backward-thread",
+        action="store_true",
+        help="Run loss.backward() on a worker thread",
+    )
+    args = parser.parse_args(argv)
+    if args.user_range and roctx is None:
+        parser.error("--user-range requires ROCTX Python")
+    return args
+
+
+def run_backward(loss, caught):
+    """Thread target: run loss.backward() and keep any exception in caught."""
+    try:
+        loss.backward()
+    except BaseException:
+        caught.append(sys.exc_info()[1:])
+
+
+def backward_on_worker_thread(loss):
+    """Run loss.backward() on a worker thread and re-raise its exception here."""
+    caught = []
+    worker = threading.Thread(target=run_backward, args=(loss, caught))
+    worker.start()
+    worker.join()
+    if caught:
+        err, tb = caught[0]
+        raise err.with_traceback(tb)
+
+
+def main(argv=None):
+    args = parse_args(argv)
+    if not torch.cuda.is_available():
+        print("GPU is required for this sample. Exiting.")
+        sys.exit(1)
+
+    if args.user_range:
+        roctx.rangePush("TrainingLoop_UserDefinedMarker")
+
+    try:
+        model = SimpleNet(same_line_linear=args.same_line_linear).cuda()
+        x = torch.randn(5, 10).cuda()
+
+        for _ in range(1):
+            output = model(x)
+            loss = output.sum()
+            if args.backward_thread:
+                backward_on_worker_thread(loss)
+            elif not args.no_backward:
+                loss.backward()
+    finally:
+        if args.user_range:
+            roctx.rangePop()
+
+    print("Training completed")
+
+
+if __name__ == "__main__":
+    main()
diff --git a/skills/rocprof-compute-sample/vcopy.cpp b/skills/rocprof-compute-sample/vcopy.cpp
new file mode 100644
index 00000000..a5109ba3
--- /dev/null
+++ b/skills/rocprof-compute-sample/vcopy.cpp
@@ -0,0 +1,227 @@
+// Copyright (c) Advanced Micro Devices, Inc.
+// SPDX-License-Identifier:  MIT
+
+#include "hip/hip_runtime.h"
+#include <assert.h>
+#include <thread>
+#include <stdio.h>
+#include <stdlib.h>
+#include <math.h>
+#include <iostream>
+using namespace std;
+
+#define HIP_ASSERT(x) (assert((x)==hipSuccess))
+
+// Unlike HIP_ASSERT, this reports the failing call and error string, and stays
+// active when NDEBUG is defined.
+#define HIP_CHECK(x)                                                           \
+  do {                                                                         \
+    hipError_t _hip_check_err = (x);                                           \
+    if (_hip_check_err != hipSuccess) {                                        \
+      fflush(stdout); /* keep the error in sequence when stdout is a pipe */   \
+      fprintf(stderr, "HIP error at %s:%d in '%s': %s (%d)\n", __FILE__,       \
+              __LINE__, #x, hipGetErrorString(_hip_check_err),                 \
+              (int)_hip_check_err);                                            \
+      exit(EXIT_FAILURE);                                                      \
+    }                                                                          \
+  } while (0)
+
+// Cap on how many individual mismatches are printed. A failed launch leaves
+// every element mismatched, which is millions of lines for a typical -n.
+#define MAX_REPORTED_ERRORS 10
+
+// HIP kernel. Each thread takes care of one element of c
+__global__ void vecCopy(double *a, double *b, double *c, int n, int stride) {
+    // Get our global thread ID
+    int id = blockIdx.x*blockDim.x+threadIdx.x;
+    if (id < n)
+      c[id] = a[id];
+}
+
+// Duplicate of vecCopy kernel. Included for testing purposes
+__global__ void vecCopy_2(double *a, double *b, double *c, int n, int stride) {
+    // Get our global thread ID
+    int id = blockIdx.x*blockDim.x+threadIdx.x;
+    if (id < n)
+      c[id] = a[id];
+}
+
+void usage() {
+  std::cout << "Usage: vcopy [OPTIONS]\n";
+  std::cout << "Required:\n";
+  std::cout << "  -n/--numThreads <value>   Set the num of threads\n";
+  std::cout << "  -b/--blockSize <value>    Set the block size\n";
+  std::cout << "Optional:\n";
+  std::cout << "  -d/--dev <value>          Set the device ID [Default: 0]\n";
+  std::cout << "  -i/--iter <value>         Set the num of iterations [Default: 1]\n";
+  std::cout << "  -h/--help                 Display this help message\n";
+  exit(1);
+  return;
+}
+
+int main(int argc, char* argv[]) {
+  // Size of vectors
+  int n; //64 MB
+  int blockSize, gridSize;
+
+  // Launch multiple kernels
+  bool multiKernel = false;
+
+  // Host input vectors
+  double *h_a;
+  double *h_b;
+  //Host output vector
+  double *h_c;
+  //Host output vector for verification
+  double *h_verify_c;
+
+  // Device input vectors
+  double *d_a;
+  double *d_b;
+  // Device output vector
+  double *d_c;
+
+  int stride = 1;
+  int devId = 0;
+  int numIter = 1;
+
+  for (int i = 0; i < argc; i++){
+    std::string arg = argv[i];
+    if ((arg == "--blockSize" || arg == "-b") && i+1 < argc)
+      blockSize = std::atoi(argv[i+1]);
+
+    else if ((arg == "--vec" || arg == "-n") && i+1 < argc)
+      n = std::atoi(argv[i+1]);
+
+    else if ((arg == "--device" || arg == "-d") && i+1 < argc)
+      devId = std::atoi(argv[i+1]);
+
+    else if ((arg == "--iter" || arg == "-i") && i+1 < argc)
+      numIter = std::atoi(argv[i+1]);
+
+    else if (arg == "--multikernel")
+      multiKernel = true;
+
+    else if (arg == "--help" || arg == "-h")
+      usage();
+  }
+
+  if (blockSize == 0)
+    usage();
+
+  if (n == 0)
+    usage();
+
+
+  int numGpuDevices;
+  HIP_ASSERT(hipGetDeviceCount(&numGpuDevices));
+  if(devId >= numGpuDevices)
+    devId = 0;
+  HIP_ASSERT(hipSetDevice(devId));
+
+  printf("vcopy testing on GCD %d\n", devId);
+
+  assert(n > 0);
+  assert(blockSize > 0);
+
+  // Size, in bytes, of each vector
+  size_t bytes = n*sizeof(double)*stride;
+
+  // Allocate memory for each vector on host
+  h_a = (double*)malloc(bytes);
+  h_b = (double*)malloc(bytes);
+  h_c = (double*)malloc(bytes);
+  h_verify_c = (double*)malloc(bytes);
+
+  printf("Finished allocating vectors on the CPU\n");
+
+  // Allocate memory for each vector on GPU
+  HIP_ASSERT(hipMalloc(&d_a, bytes));
+  HIP_ASSERT(hipMalloc(&d_b, bytes));
+  HIP_ASSERT(hipMalloc(&d_c, bytes));
+
+  printf("Finished allocating vectors on the GPU\n");
+
+  // Initialize vectors on host
+  for(int i = 0; i < n; i++) {
+      h_a[i] = i;
+      h_b[i] = i;
+  }
+
+  // Copy host vectors to device
+  HIP_ASSERT(hipMemcpy(d_a, h_a, bytes, hipMemcpyHostToDevice));
+  HIP_ASSERT(hipMemcpy(d_b, h_b, bytes, hipMemcpyHostToDevice));
+
+  printf("Finished copying vectors to the GPU\n");
+
+  // Number of thread blocks in grid
+  gridSize = (int)ceil((float)n/blockSize);
+  int tot_waves = (blockSize*gridSize)/64;
+  float num_bytes_kb = ((sizeof(double))*n)/(1024);
+  float num_bytes_wave = (1.0*num_bytes_kb)/(1.0*tot_waves);
+
+  printf("sw thinks it moved %f KB per wave \n", (2.0*num_bytes_wave));
+  printf("Total threads: %d, Grid Size: %d block Size:%d, Wavefronts:%d:\n", n, gridSize, blockSize, tot_waves);
+  printf("Launching the  kernel on the GPU\n");
+
+  // Execute the kernel
+  for(int i = 0; i < numIter; i++){
+    hipLaunchKernelGGL(vecCopy, dim3(gridSize), dim3(blockSize), 0, 0, d_a, d_b, d_c, n, stride);
+    // Catches launch-time failures (e.g. no kernel image for this device);
+    // hipDeviceSynchronize() catches failures during execution.
+    HIP_CHECK(hipGetLastError());
+    HIP_CHECK(hipDeviceSynchronize());
+    printf("Finished executing kernel\n");
+    // Optionally, launch a second kernel. Only here for testing purposes
+    if (multiKernel){
+      hipLaunchKernelGGL(vecCopy_2, dim3(gridSize), dim3(blockSize), 0, 0, d_a, d_b, d_c, n, stride);
+      HIP_CHECK(hipGetLastError());
+      HIP_CHECK(hipDeviceSynchronize());
+      printf("Finished executing kernel\n");
+    }
+  }
+
+  // Copy array back to host
+  HIP_ASSERT(hipMemcpy( h_c, d_c, bytes, hipMemcpyDeviceToHost));
+  printf("Finished copying the output vector from the GPU to the CPU\n");
+
+  // Compute for CPU
+  for(int i=0; i<n; i++) {
+    // h_verify_c[i*stride] = h_a[i*stride] + h_b[i*stride];
+    h_verify_c[i*stride] = h_a[i*stride] ;
+  }
+
+  // Verfiy results
+  int numErrors = 0;
+  for(int i = 0; i < n; i++) {
+    if (abs(h_verify_c[i*stride] - h_c[i*stride]) > 1e-5) {
+      if (numErrors < MAX_REPORTED_ERRORS)
+        printf("Error at position i %d, Expected: %f, Found: %f \n", i,
+               h_verify_c[i*stride], h_c[i*stride]);
+      numErrors++;
+    }
+  }
+  if (numErrors > MAX_REPORTED_ERRORS)
+    printf("... %d further mismatches not shown\n", numErrors - MAX_REPORTED_ERRORS);
+  if (numErrors > 0)
+    printf("Verification FAILED: %d of %d elements mismatched\n", numErrors, n);
+  //printf("Printing few elements from the output vector\n");
+  for(int i = 0; i < 20; i++) {
+    //printf("Output[%d]:%f\n",i, h_c[i]);
+  }
+
+  printf("Releasing GPU memory\n");
+
+  // Release device memory
+  HIP_ASSERT(hipFree(d_a));
+  HIP_ASSERT(hipFree(d_b));
+  HIP_ASSERT(hipFree(d_c));
+
+  // Release host memory
+  printf("Releasing CPU memory\n");
+  free(h_a);
+  free(h_b);
+  free(h_c);
+
+  return numErrors > 0 ? EXIT_FAILURE : EXIT_SUCCESS;
+}
diff --git a/skills/rocprof-compute-sample/vmem.hip b/skills/rocprof-compute-sample/vmem.hip
new file mode 100644
index 00000000..4f079cf7
--- /dev/null
+++ b/skills/rocprof-compute-sample/vmem.hip
@@ -0,0 +1,75 @@
+// Copyright (c) Advanced Micro Devices, Inc.
+// SPDX-License-Identifier:  MIT
+
+/*
+An example to explore global/generic instructions.
+Written by Nicholas Curtis [AMD].
+*/
+
+#include "common.h"
+
+typedef int __attribute__((address_space(0)))* generic_ptr;
+
+__attribute__((noinline)) __device__ void generic_store(generic_ptr ptr, int zero) { *ptr = zero; }
+__attribute__((noinline)) __device__ int generic_load(generic_ptr ptr) { return *ptr; }
+__attribute__((noinline)) __device__ void generic_atomic(generic_ptr ptr, int zero) { atomicAdd((int*)ptr, zero); }
+
+__global__ void global_write(int* ptr, int zero) {
+  ptr[threadIdx.x] = zero;
+}
+
+__global__ void generic_write(int* ptr, int zero, int filter) {
+  __shared__ int lds[1024];
+  int* generic = (threadIdx.x < filter) ? &ptr[threadIdx.x] : &lds[threadIdx.x];
+  generic_store((generic_ptr)generic, zero);
+}
+
+__global__ void global_read(int* ptr, int zero) {
+  int x = ptr[threadIdx.x];
+  if (x != zero) {
+    ptr[threadIdx.x] = x + 1;
+  }
+}
+
+__global__ void generic_read(int* ptr, int zero, int filter) {
+  __shared__ int lds[1024];
+  if (static_cast<int>(filter - 1) == zero) {
+    lds[threadIdx.x] = 0; // initialize to zero to avoid conditional, but hide behind _another_ conditional
+  }
+  int* generic;
+  if (static_cast<int>(threadIdx.x) > filter - 1) {
+    generic = &ptr[threadIdx.x];
+  } else {
+    generic = &lds[threadIdx.x];
+    abort();
+  }
+  int x = generic_load((generic_ptr)generic);
+  if (x != zero) {
+    ptr[threadIdx.x] = x + 1;
+  }
+}
+
+
+__global__ void global_atomic(int* ptr, int zero) {
+  atomicAdd(ptr, zero);
+}
+
+__global__ void generic_atomic(int* ptr, int filter, int zero) {
+  __shared__ int lds[1024];
+  int* generic = (threadIdx.x % 2 == filter) ? &ptr[threadIdx.x] : &lds[threadIdx.x];
+  generic_atomic((generic_ptr)generic, zero);
+}
+
+int main() {
+  int* ptr;
+  hipCheck(hipMalloc(&ptr, sizeof(int)));
+  hipCheck(hipMemset(ptr, 0, sizeof(int)));
+  global_write<<<1,1>>>(ptr, 0);
+  generic_write<<<1,1>>>(ptr, 0, 0);
+  global_read<<<1,1>>>(ptr, 0);
+  generic_read<<<1,1>>>(ptr, 0, 0);
+  global_atomic<<<1,1>>>(ptr, 0);
+  generic_atomic<<<1,1>>>(ptr, 0, 0);
+  hipCheck(hipDeviceSynchronize());
+  hipCheck(hipFree(ptr));
+}
diff --git a/skills/rocprof-compute-sample/vrandom_access.cpp b/skills/rocprof-compute-sample/vrandom_access.cpp
new file mode 100644
index 00000000..c53f54b2
--- /dev/null
+++ b/skills/rocprof-compute-sample/vrandom_access.cpp
@@ -0,0 +1,82 @@
+// Copyright (c) Advanced Micro Devices, Inc.
+// SPDX-License-Identifier:  MIT
+
+/*
+An example code to execute random access to explore cache hits/misses in L2 Cache.
+*/
+
+#include <hip/hip_runtime.h>
+#include <assert.h>
+#include <iostream>
+#include <cstdlib>
+#include <ctime>
+#include <vector>
+#include <algorithm>
+
+#define HIP_ASSERT(x) (assert((x) == hipSuccess))
+
+// Kernel: random access, each thread picks a random index
+__global__ void randomAccessKernel(int *d_data, int N, unsigned int *d_seeds)
+{
+    int tid = blockIdx.x * blockDim.x + threadIdx.x;
+    if (tid < N)
+    {
+        unsigned int seed = d_seeds[tid];
+        // Simple XORShift
+        seed ^= (seed << 13);
+        seed ^= (seed >> 17);
+        seed ^= (seed << 5);
+        int idx = seed % N;
+        d_data[idx] += 1;
+    }
+}
+
+int main()
+{
+    const int N = 1 << 24; // Try 16M elements to exceed cache
+    size_t size = N * sizeof(int);
+
+    // Host memory
+    std::vector<int> h_data(N, 0);
+    std::vector<unsigned int> h_seeds(N);
+
+    // Generate seeds
+    srand(time(nullptr));
+    for (int i = 0; i < N; ++i)
+    {
+        // Keep them diverse. Could be random or based on i
+        h_seeds[i] = rand();
+    }
+
+    // Allocate device memory
+    int *d_data;
+    unsigned int *d_seeds;
+    HIP_ASSERT(hipMalloc(&d_data, size));
+    HIP_ASSERT(hipMalloc(&d_seeds, N * sizeof(unsigned int)));
+
+    // Copy h_data to device
+    HIP_ASSERT(hipMemcpy(d_data, h_data.data(), size, hipMemcpyHostToDevice));
+    HIP_ASSERT(hipMemcpy(d_seeds, h_seeds.data(), N * sizeof(unsigned int), hipMemcpyHostToDevice));
+
+    // Configure kernel
+    dim3 blockSize(64);
+    dim3 gridSize((N + blockSize.x - 1) / blockSize.x);
+
+    // Repeat the launch so the profile is robust.
+    const int kIters = 30;
+    for (int i = 0; i < kIters; ++i)
+    {
+        hipLaunchKernelGGL(randomAccessKernel, gridSize, blockSize, 0, 0, d_data, N, d_seeds);
+    }
+    HIP_ASSERT(hipDeviceSynchronize());
+
+    HIP_ASSERT(hipMemcpy(h_data.data(), d_data, size, hipMemcpyDeviceToHost));
+
+    // Cleanup
+    HIP_ASSERT(hipFree(d_data));
+    HIP_ASSERT(hipFree(d_seeds));
+    ;
+
+    std::cout << "RandomAccess HIP test completed.\n";
+    return 0;
+}
diff --git a/skills/rocprof-compute-sample/vsequential_access.cpp b/skills/rocprof-compute-sample/vsequential_access.cpp
new file mode 100644
index 00000000..d443fc30
--- /dev/null
+++ b/skills/rocprof-compute-sample/vsequential_access.cpp
@@ -0,0 +1,60 @@
+// Copyright (c) Advanced Micro Devices, Inc.
+// SPDX-License-Identifier:  MIT
+
+/*
+An example code to execute sequential access to explore cache hits/misses in L2 Cache.
+*/
+
+#include <hip/hip_runtime.h>
+#include <iostream>
+#include <assert.h>
+
+#define HIP_ASSERT(x) (assert((x)==hipSuccess))
+
+// Kernel: sequential access, each thread reads/writes an element in order
+__global__ void sequentialAccessKernel(int *d_data, int N)
+{
+    int tid = blockIdx.x * blockDim.x + threadIdx.x;
+    if (tid < N)
+    {
+        d_data[tid] += 1;
+    }
+}
+
+int main()
+{
+    const int N = 1 << 20; // 1M elements
+    size_t size = N * sizeof(int);
+    // Allocate host memory
+    int *h_data = (int *)malloc(size);
+    std::fill_n(h_data, N, 0);
+
+    // Allocate device memory
+    int *d_data;
+    HIP_ASSERT(hipMalloc(&d_data, size));
+
+    // Copy h_data to device
+    HIP_ASSERT(hipMemcpy(d_data, h_data, size, hipMemcpyHostToDevice));
+
+    // Configure kernel
+    dim3 blockSize(64);
+    dim3 gridSize((N + blockSize.x - 1) / blockSize.x);
+
+    // Repeat the launch so the profile is robust.
+    const int kIters = 30;
+    for (int i = 0; i < kIters; ++i)
+    {
+        hipLaunchKernelGGL(sequentialAccessKernel, gridSize, blockSize, 0, 0, d_data, N);
+    }
+    HIP_ASSERT(hipDeviceSynchronize());
+
+    // Copy back to host
+    HIP_ASSERT(hipMemcpy(h_data, d_data, size, hipMemcpyDeviceToHost));
+
+    // Cleanup
+    HIP_ASSERT(hipFree(d_data));
+    free(h_data);
+
+    std::cout << "SequentialAccess HIP test completed.\n";
+    return 0;
+}
diff --git a/skills/rocprof-compute-sample/wmma_gemm_fp16/Makefile b/skills/rocprof-compute-sample/wmma_gemm_fp16/Makefile
new file mode 100644
index 00000000..4266c564
--- /dev/null
+++ b/skills/rocprof-compute-sample/wmma_gemm_fp16/Makefile
@@ -0,0 +1,34 @@
+# Copyright (c) Advanced Micro Devices, Inc.
+# SPDX-License-Identifier: MIT
+#
+# FP16 WMMA GEMM sample for gfx1250 metric health profiling.
+
+HIPCC      ?= hipcc
+GPU_ARCH   ?= gfx1250
+HIPCC_FLAGS = -O2 -g -std=c++17 -Wall -Wextra -fPIC
+INCLUDES    = -I.
+
+TARGET      = wmma_gemm_fp16
+KERNEL_SRC  = wmma_gemm_fp16.hip
+HOST_SRC    = main.cpp
+
+.PHONY: all clean run help
+
+all: $(TARGET)
+
+$(TARGET): $(KERNEL_SRC) $(HOST_SRC) wmma_gemm_fp16_config.hpp wmma_gemm_fp16.hpp
+	@echo "Building $(TARGET) for $(GPU_ARCH)..."
+	$(HIPCC) $(HIPCC_FLAGS) --offload-arch=$(GPU_ARCH) $(INCLUDES) -o $@ $(HOST_SRC) $(KERNEL_SRC)
+	@echo "Build complete: $@"
+
+run: $(TARGET)
+	./$(TARGET)
+
+clean:
+	rm -f $(TARGET)
+
+help:
+	@echo "Targets:"
+	@echo "  make            Build for gfx1250 (override: GPU_ARCH=gfx1250)"
+	@echo "  make run        Build and run with defaults"
+	@echo "  make clean      Remove binary"
diff --git a/skills/rocprof-compute-sample/wmma_gemm_fp16/main.cpp b/skills/rocprof-compute-sample/wmma_gemm_fp16/main.cpp
new file mode 100644
index 00000000..c6175364
--- /dev/null
+++ b/skills/rocprof-compute-sample/wmma_gemm_fp16/main.cpp
@@ -0,0 +1,208 @@
+// Copyright (c) Advanced Micro Devices, Inc.
+// SPDX-License-Identifier:  MIT
+
+#include <chrono>
+#include <cmath>
+#include <cstdlib>
+#include <iostream>
+#include <random>
+#include <string>
+#include <vector>
+
+#include "../common.h"
+#include "wmma_gemm_fp16.hpp"
+
+using namespace wmma_gemm_fp16;
+
+namespace {
+
+struct Options {
+    uint32_t m     = kDefaultM;
+    uint32_t n     = kDefaultN;
+    uint32_t k     = kDefaultK;
+    uint32_t iters = kDefaultIters;
+    uint32_t warmup = kDefaultWarmup;
+    int device   = 0;
+};
+
+void usage()
+{
+    std::cout
+        << "Usage: wmma_gemm_fp16 [OPTIONS]\n"
+        << "  FP16 WMMA GEMM for gfx1250 WMMA/XDL metric health profiling.\n\n"
+        << "Optional:\n"
+        << "  -m, --m <value>       Rows of A / D [default: " << kDefaultM << "]\n"
+        << "  -n, --n <value>       Cols of B / D [default: " << kDefaultN << "]\n"
+        << "  -k, --k <value>       Inner dimension K (multiple of 32) [default: "
+        << kDefaultK << "]\n"
+        << "  -i, --iter <value>    Kernel iterations [default: " << kDefaultIters
+        << "]\n"
+        << "  -w, --warmup <value>  Warmup iterations [default: " << kDefaultWarmup
+        << "]\n"
+        << "  -d, --device <id>     HIP device id [default: 0]\n"
+        << "  -h, --help            Show this help\n";
+    std::exit(EXIT_SUCCESS);
+}
+
+Options parse_args(int argc, char* argv[])
+{
+    Options opts;
+    for (int i = 1; i < argc; ++i) {
+        const std::string arg = argv[i];
+        auto need_value       = [&](const char* flag) {
+            if (i + 1 >= argc) {
+                std::cerr << "Missing value for " << flag << "\n";
+                usage();
+            }
+            return std::string(argv[++i]);
+        };
+
+        if (arg == "-h" || arg == "--help") {
+            usage();
+        } else if (arg == "-m" || arg == "--m") {
+            opts.m = static_cast<uint32_t>(std::stoul(need_value(arg.c_str())));
+        } else if (arg == "-n" || arg == "--n") {
+            opts.n = static_cast<uint32_t>(std::stoul(need_value(arg.c_str())));
+        } else if (arg == "-k" || arg == "--k") {
+            opts.k = static_cast<uint32_t>(std::stoul(need_value(arg.c_str())));
+        } else if (arg == "-i" || arg == "--iter") {
+            opts.iters = static_cast<uint32_t>(std::stoul(need_value(arg.c_str())));
+        } else if (arg == "-w" || arg == "--warmup") {
+            opts.warmup = static_cast<uint32_t>(std::stoul(need_value(arg.c_str())));
+        } else if (arg == "-d" || arg == "--device") {
+            opts.device = std::stoi(need_value(arg.c_str()));
+        } else {
+            std::cerr << "Unknown argument: " << arg << "\n";
+            usage();
+        }
+    }
+    return opts;
+}
+
+void fill_matrix(std::vector<_Float16>& mat, uint32_t rows, uint32_t cols)
+{
+    std::mt19937 rng(1337);
+    std::uniform_real_distribution<float> dist(0.5f, 2.0f);
+    mat.resize(static_cast<size_t>(rows) * cols);
+    for (auto& value : mat) {
+        value = static_cast<_Float16>(dist(rng));
+    }
+}
+
+bool smoke_check(const std::vector<_Float16>& d, uint32_t m, uint32_t n)
+{
+    for (uint32_t row = 0; row < std::min(m, 64u); ++row) {
+        for (uint32_t col = 0; col < std::min(n, 64u); ++col) {
+            const float value = static_cast<float>(d[static_cast<size_t>(row) * n + col]);
+            if (value != 0.0f && std::isfinite(value)) {
+                return true;
+            }
+        }
+    }
+    return false;
+}
+
+}  // namespace
+
+int main(int argc, char* argv[])
+{
+    const Options opts = parse_args(argc, argv);
+
+    if (opts.m == 0 || opts.n == 0 || opts.k == 0) {
+        std::cerr << "M, N, and K must be positive.\n";
+        return EXIT_FAILURE;
+    }
+    if (opts.k % kBlockK != 0) {
+        std::cerr << "K must be a multiple of " << kBlockK << " for 16x16x32 WMMA tiles.\n";
+        return EXIT_FAILURE;
+    }
+    if (opts.m % kBlockM != 0 || opts.n % kBlockN != 0) {
+        std::cerr << "M and N must be multiples of " << kBlockM << ".\n";
+        return EXIT_FAILURE;
+    }
+
+    hipCheck(hipSetDevice(opts.device));
+
+    hipDeviceProp_t props{};
+    hipCheck(hipGetDeviceProperties(&props, opts.device));
+    std::cout << "Device: " << props.name << " (" << props.gcnArchName << ")\n";
+    std::cout << "Problem: M=" << opts.m << " N=" << opts.n << " K=" << opts.k
+              << " iters=" << opts.iters << "\n";
+
+    std::vector<_Float16> h_a;
+    std::vector<_Float16> h_b;
+    std::vector<_Float16> h_d(static_cast<size_t>(opts.m) * opts.n, _Float16(0));
+
+    fill_matrix(h_a, opts.m, opts.k);
+    fill_matrix(h_b, opts.n, opts.k);
+
+    _Float16 *d_a = nullptr;
+    _Float16 *d_b = nullptr;
+    _Float16 *d_d = nullptr;
+
+    const size_t bytes_a = h_a.size() * sizeof(_Float16);
+    const size_t bytes_b = h_b.size() * sizeof(_Float16);
+    const size_t bytes_d = h_d.size() * sizeof(_Float16);
+
+    hipCheck(hipMalloc(&d_a, bytes_a));
+    hipCheck(hipMalloc(&d_b, bytes_b));
+    hipCheck(hipMalloc(&d_d, bytes_d));
+
+    hipCheck(hipMemcpy(d_a, h_a.data(), bytes_a, hipMemcpyHostToDevice));
+    hipCheck(hipMemcpy(d_b, h_b.data(), bytes_b, hipMemcpyHostToDevice));
+    hipCheck(hipMemcpy(d_d, h_d.data(), bytes_d, hipMemcpyHostToDevice));
+
+    const dim3 block_dim(kThreadsX, kThreadsY);
+    const dim3 grid_dim(DivUp(static_cast<int>(opts.m), static_cast<int>(kBlockM)),
+                        DivUp(static_cast<int>(opts.n), static_cast<int>(kBlockN)),
+                        1);
+
+    std::cout << "Launch grid=(" << grid_dim.x << "," << grid_dim.y << ") block=("
+              << block_dim.x << "," << block_dim.y << ")\n";
+
+    auto launch = [&]() {
+        hipLaunchKernelGGL(gemm_wmma_fp16,
+                           grid_dim,
+                           block_dim,
+                           0,
+                           0,
+                           opts.m,
+                           opts.n,
+                           opts.k,
+                           d_a,
+                           d_b,
+                           d_d);
+        hipCheck(hipGetLastError());
+    };
+
+    for (uint32_t i = 0; i < opts.warmup; ++i) {
+        launch();
+    }
+    hipCheck(hipDeviceSynchronize());
+
+    const auto start = std::chrono::steady_clock::now();
+    for (uint32_t i = 0; i < opts.iters; ++i) {
+        launch();
+    }
+    hipCheck(hipDeviceSynchronize());
+    const auto end = std::chrono::steady_clock::now();
+
+    const double elapsed_ms =
+        std::chrono::duration<double, std::milli>(end - start).count();
+    const double gflops =
+        2.0 * static_cast<double>(opts.m) * opts.n * opts.k * opts.iters * 1e-9;
+    const double tflops_per_s = gflops / elapsed_ms;
+
+    hipCheck(hipMemcpy(h_d.data(), d_d, bytes_d, hipMemcpyDeviceToHost));
+
+    const bool passed = smoke_check(h_d, opts.m, opts.n);
+    std::cout << "Elapsed: " << elapsed_ms << " ms (" << opts.iters << " iterations)\n";
+    std::cout << "Throughput: " << tflops_per_s << " TFLOP/s (FP16 GEMM convention)\n";
+    std::cout << "Smoke check: " << (passed ? "PASS (non-zero output)" : "FAIL") << "\n";
+
+    hipCheck(hipFree(d_a));
+    hipCheck(hipFree(d_b));
+    hipCheck(hipFree(d_d));
+
+    return passed ? EXIT_SUCCESS : EXIT_FAILURE;
+}
diff --git a/skills/rocprof-compute-sample/wmma_gemm_fp16/wmma_gemm_fp16.hip b/skills/rocprof-compute-sample/wmma_gemm_fp16/wmma_gemm_fp16.hip
new file mode 100644
index 00000000..eb6ef300
--- /dev/null
+++ b/skills/rocprof-compute-sample/wmma_gemm_fp16/wmma_gemm_fp16.hip
@@ -0,0 +1,153 @@
+// Copyright (c) Advanced Micro Devices, Inc.
+// SPDX-License-Identifier:  MIT
+
+// Sustained FP16 WMMA GEMM for gfx1250 metric health profiling.
+// Adapted from hipNPIKernels WMMA_gemm fp16_16x16x32 (native16x16x32).
+
+#include <hip/hip_fp16.h>
+#include <hip/hip_runtime.h>
+
+#include "wmma_gemm_fp16_config.hpp"
+
+using namespace wmma_gemm_fp16;
+
+using Float16x4  = _Float16 __attribute__((ext_vector_type(4)));
+using Float16x8  = _Float16 __attribute__((ext_vector_type(8)));
+using Float16x16 = _Float16 __attribute__((ext_vector_type(16)));
+using Floatx8    = float __attribute__((ext_vector_type(8)));
+
+__device__ inline uint32_t lane_id() { return static_cast<uint32_t>(threadIdx.x); }
+
+__device__ inline void load_row_major(Float16x16& regs,
+                                      uint32_t row,
+                                      uint32_t k_start,
+                                      uint32_t row_len,
+                                      const _Float16* __restrict__ buffer)
+{
+    const uint32_t low_high = lane_id() / 16;
+    const uint32_t row_addr = row + (lane_id() % 16);
+
+    const uint32_t col1 = k_start + low_high * 16;
+    const uint32_t col2 = k_start + 4 + low_high * 16;
+    const uint32_t col3 = k_start + 8 + low_high * 16;
+    const uint32_t col4 = k_start + 12 + low_high * 16;
+
+    auto* p_regs = reinterpret_cast<Float16x4*>(&regs);
+    p_regs[0]    = *reinterpret_cast<const Float16x4*>(buffer + row_addr * row_len + col1);
+    p_regs[1]    = *reinterpret_cast<const Float16x4*>(buffer + row_addr * row_len + col2);
+    p_regs[2]    = *reinterpret_cast<const Float16x4*>(buffer + row_addr * row_len + col3);
+    p_regs[3]    = *reinterpret_cast<const Float16x4*>(buffer + row_addr * row_len + col4);
+}
+
+__device__ inline void write_output(const Float16x8& frag,
+                                    uint32_t wave_row,
+                                    uint32_t wave_col,
+                                    uint32_t row_len,
+                                    _Float16* __restrict__ d)
+{
+    const uint32_t out_col = (lane_id() >= 16) ? 8 : 0;
+    const uint32_t out_row = lane_id() % 16;
+    const uint32_t row     = wave_row + out_row;
+    const uint32_t col     = wave_col + out_col;
+    *reinterpret_cast<Float16x8*>(d + row * row_len + col) = frag;
+}
+
+__device__ inline void inner_gemm(uint32_t sub_m_start,
+                                  uint32_t sub_n_start,
+                                  uint32_t m,
+                                  uint32_t n,
+                                  uint32_t k,
+                                  const _Float16* __restrict__ a,
+                                  const _Float16* __restrict__ b,
+                                  _Float16* __restrict__ d)
+{
+#if !__has_builtin(__builtin_amdgcn_wmma_f32_16x16x32_f16)
+    (void)sub_m_start;
+    (void)sub_n_start;
+    (void)m;
+    (void)n;
+    (void)k;
+    (void)a;
+    (void)b;
+    if (threadIdx.x == 0 && blockIdx.x == 0 && blockIdx.y == 0) {
+        d[0] = _Float16(0);
+    }
+    return;
+#else
+    Floatx8 frag_acc = {};
+
+    const uint32_t major_wave = __builtin_amdgcn_readfirstlane(threadIdx.x / kWaveSize);
+    const uint32_t minor_wave = threadIdx.y;
+
+    const uint32_t wave_row = major_wave * kWmmaM + sub_m_start;
+    const uint32_t wave_col = minor_wave * kWmmaN + sub_n_start;
+
+    const bool full_block =
+        (sub_m_start + kBlockM <= m) && (sub_n_start + kBlockN <= n) && (k >= kBlockK);
+
+    if (!full_block) {
+        return;
+    }
+
+    const uint32_t k_tiles = k / kBlockK;
+    uint32_t k_iter        = 0;
+
+    Float16x16 reg_a = {};
+    Float16x16 reg_b = {};
+
+    load_row_major(reg_a, wave_row, 0, k, a);
+    load_row_major(reg_b, wave_col, 0, k, b);
+
+    if (k_tiles > 1) {
+        do {
+            const uint32_t next_k = (k_iter + 1) * kBlockK;
+
+            Float16x16 next_a = {};
+            Float16x16 next_b = {};
+
+            load_row_major(next_a, wave_row, next_k, k, a);
+            load_row_major(next_b, wave_col, next_k, k, b);
+
+            __builtin_amdgcn_sched_barrier(0);
+            frag_acc = __builtin_amdgcn_wmma_f32_16x16x32_f16(
+                0, reg_b, 0, reg_a, 0, frag_acc, 0, 0);
+            __builtin_amdgcn_sched_barrier(0);
+
+            reg_a = next_a;
+            reg_b = next_b;
+            ++k_iter;
+        } while (k_iter < (k_tiles - 1));
+    }
+
+    frag_acc = __builtin_amdgcn_wmma_f32_16x16x32_f16(0, reg_b, 0, reg_a, 0, frag_acc, 0, 0);
+
+    Float16x8 frag_d = {};
+    for (int i = 0; i < 8; ++i) {
+        frag_d[i] = static_cast<_Float16>(frag_acc[i]);
+    }
+
+    write_output(frag_d, wave_row, wave_col, n, d);
+#endif
+}
+
+__attribute__((amdgpu_flat_work_group_size(1, kFlatWorkgroupSize)))
+extern "C" __global__ void gemm_wmma_fp16(uint32_t m,
+                                          uint32_t n,
+                                          uint32_t k,
+                                          const _Float16* __restrict__ a,
+                                          const _Float16* __restrict__ b,
+                                          _Float16* __restrict__ d)
+{
+    const uint32_t block_count_m = DivUp(static_cast<int>(m), static_cast<int>(kBlockM));
+    const uint32_t block_count_n = DivUp(static_cast<int>(n), static_cast<int>(kBlockN));
+
+    for (uint32_t bm = blockIdx.x; bm < block_count_m; bm += gridDim.x) {
+        for (uint32_t bn = blockIdx.y; bn < block_count_n; bn += gridDim.y) {
+            const uint32_t block_m = bm * kBlockM;
+            const uint32_t block_n = bn * kBlockN;
+            if (block_m < m && block_n < n) {
+                inner_gemm(block_m, block_n, m, n, k, a, b, d);
+            }
+        }
+    }
+}
diff --git a/skills/rocprof-compute-sample/wmma_gemm_fp16/wmma_gemm_fp16.hpp b/skills/rocprof-compute-sample/wmma_gemm_fp16/wmma_gemm_fp16.hpp
new file mode 100644
index 00000000..faacb53b
--- /dev/null
+++ b/skills/rocprof-compute-sample/wmma_gemm_fp16/wmma_gemm_fp16.hpp
@@ -0,0 +1,19 @@
+// Copyright (c) Advanced Micro Devices, Inc.
+// SPDX-License-Identifier: MIT
+
+#pragma once
+
+#include <cstdint>
+
+#include <hip/hip_runtime.h>
+
+#include "wmma_gemm_fp16_config.hpp"
+
+// FP16 WMMA GEMM for gfx1250. Kernel adapted from hipNPIKernels
+// WMMA_gemm/src/fp16/gfx1250/fp16_16x16x32/native16x16x32.cpp.
+extern "C" __global__ void gemm_wmma_fp16(uint32_t m,
+                                          uint32_t n,
+                                          uint32_t k,
+                                          const _Float16* __restrict__ a,
+                                          const _Float16* __restrict__ b,
+                                          _Float16* __restrict__ d);
diff --git a/skills/rocprof-compute-sample/wmma_gemm_fp16/wmma_gemm_fp16_config.hpp b/skills/rocprof-compute-sample/wmma_gemm_fp16/wmma_gemm_fp16_config.hpp
new file mode 100644
index 00000000..2bb316e6
--- /dev/null
+++ b/skills/rocprof-compute-sample/wmma_gemm_fp16/wmma_gemm_fp16_config.hpp
@@ -0,0 +1,31 @@
+// Copyright (c) Advanced Micro Devices, Inc.
+// SPDX-License-Identifier: MIT
+
+#pragma once
+
+#include <cstdint>
+
+namespace wmma_gemm_fp16 {
+
+constexpr uint32_t kWmmaM = 16;
+constexpr uint32_t kWmmaN = 16;
+constexpr uint32_t kWmmaK = 32;
+
+constexpr uint32_t kWaveSize = 32;
+constexpr uint32_t kBlockM = 16;
+constexpr uint32_t kBlockN = 16;
+constexpr uint32_t kBlockK = 32;
+
+constexpr uint32_t kThreadsX = 1 * kWaveSize;
+constexpr uint32_t kThreadsY = 1;
+constexpr uint32_t kFlatWorkgroupSize = kThreadsX * kThreadsY;
+
+constexpr uint32_t kDefaultM = 4096;
+constexpr uint32_t kDefaultN = 4096;
+constexpr uint32_t kDefaultK = 1024;
+constexpr uint32_t kDefaultIters = 50;
+constexpr uint32_t kDefaultWarmup = 5;
+
+constexpr int DivUp(int a, int b) { return (a + b - 1) / b; }
+
+}  // namespace wmma_gemm_fp16
diff --git a/skills/rocprof-compute-speed-of-light/.federated.json b/skills/rocprof-compute-speed-of-light/.federated.json
new file mode 100644
index 00000000..975c16eb
--- /dev/null
+++ b/skills/rocprof-compute-speed-of-light/.federated.json
@@ -0,0 +1,10 @@
+{
+  "source": "rocm-rocm-systems",
+  "repo": "ROCm/rocm-systems",
+  "ref": "develop",
+  "commit": "8cce7054f17a77272e1ba73ef78f4f77cf056956",
+  "path": "projects/rocprofiler-compute/skills/rocprof-compute-speed-of-light",
+  "license": "MIT",
+  "content_hash": "924898dc2f20900f8191721bf89163c4d4ac27a6fd9a758b64292657e5990818",
+  "imported_at": "2026-10-28T01:19:39Z"
+}
diff --git a/skills/rocprof-compute-speed-of-light/SKILL.md b/skills/rocprof-compute-speed-of-light/SKILL.md
new file mode 100644
index 00000000..1183abe6
--- /dev/null
+++ b/skills/rocprof-compute-speed-of-light/SKILL.md
@@ -0,0 +1,80 @@
+---
+name: rocprof-compute-speed-of-light
+description: Reads the rocprof-compute System Speed-of-Light panel to show how close an AMD GPU kernel runs to hardware peak for compute throughput, memory bandwidth, and utilization. Use when the user asks how close to peak a kernel is, what percent of peak it reaches, whether a measured percentage is good enough or normal, whether a utilization number should be higher, or asks for an overall efficiency summary. Not for CUDA tools, Windows, per-instruction analysis, or a datasheet or theoretical-peak question when there is no profile to read.
+---
+
+# Read System Speed-of-Light
+
+Speed-of-Light is the first panel to read after profiling. It puts every
+headline metric next to the hardware peak for the current GPU, so one table
+says whether a kernel is near any limit at all.
+
+This skill reads an existing workload directory. It does not collect one.
+
+Run `rocprof-compute analyze --help` before choosing flags.
+
+## 1. Get the panel
+
+```bash
+rocprof-compute analyze --path ./workloads/<name>/<gpu_model> --list-stats
+rocprof-compute analyze --path ./workloads/<name>/<gpu_model> -k <kernel_id> -b sol
+```
+
+`--path` points at the directory containing `profiling_config.yaml`. Always
+pass `-k` with one kernel id from `--list-stats`; a Speed-of-Light table
+averaged over unrelated kernels means nothing.
+
+If the panel is missing or empty, those counters were not collected. Collect
+them with the `rocprof-compute-kernel-bottleneck` skill, using `-b sol` at
+profile time for a cheap single-topic run.
+
+## 2. Read it
+
+The table has `Metric`, `Avg`, `Unit`, `Peak`, and `Percent of Peak` columns.
+The peak is already computed for the detected architecture. Compare `Avg`
+against the `Peak` in the same row and never calculate a peak yourself.
+
+`Percent of Peak` is `N/A` for metrics that are already percentages, such as
+the utilization rows. Read those as the percentage they are.
+
+What the shape of the table tells you:
+
+| Reading | Meaning | Next step |
+|---|---|---|
+| A memory row near peak | bandwidth-limited | the `rocprof-compute-memory` skill |
+| An MFMA or VALU row near peak | compute-limited | the `rocprof-compute-roofline` skill |
+| Every row far below peak | not limited by throughput | occupancy and launch limits, via the `rocprof-compute-kernel-bottleneck` skill |
+| MFMA at zero on a GEMM | matrix pipeline unused | check the kernel's instruction mix with `-b cu_ins` |
+
+A row at zero is a real reading only if the kernel was expected to use that
+unit. MFMA at zero is normal for a kernel with no matrix instructions.
+
+Do not apply remembered thresholds. "Above 70% is saturated" is not a rule the
+tool encodes, and the correct comparison is always against the printed peak,
+the neighbouring kernels, and a known-good baseline for the same workload.
+
+## 3. Understand the metrics
+
+Each metric's definition and the counters behind it are in
+[system-speed-of-light.rst](https://github.com/ROCm/rocm-systems/blob/8cce7054f17a77272e1ba73ef78f4f77cf056956/projects/rocprofiler-compute/docs/conceptual/cdna/system-speed-of-light.rst)
+for CDNA accelerators and
+[the RDNA equivalent](https://github.com/ROCm/rocm-systems/blob/8cce7054f17a77272e1ba73ef78f4f77cf056956/projects/rocprofiler-compute/docs/conceptual/rdna/system-speed-of-light.rst)
+for client APUs. The whole metric hierarchy is in the
+[performance model](https://github.com/ROCm/rocm-systems/blob/8cce7054f17a77272e1ba73ef78f4f77cf056956/projects/rocprofiler-compute/docs/conceptual/performance-model.rst).
+
+## 4. Compute Speed-of-Light
+
+The compute pipeline has its own Speed-of-Light table with the same shape:
+
+```bash
+rocprof-compute analyze --path ./workloads/<name>/<gpu_model> -k <kernel_id> -b cu_pipe
+```
+
+Read it when the system-level table shows compute near peak and you need to
+know which pipeline is responsible.
+
+## 5. Normalization
+
+`-n` changes the unit for the entire report, not one metric. Leave it at
+`per_kernel` for Speed-of-Light. The peaks are expressed to match the default,
+and switching units makes the peak comparison harder to read, not easier.
diff --git a/skills/rocprof-compute-speed-of-light/evals/evals.json b/skills/rocprof-compute-speed-of-light/evals/evals.json
new file mode 100644
index 00000000..aa24860e
--- /dev/null
+++ b/skills/rocprof-compute-speed-of-light/evals/evals.json
@@ -0,0 +1,74 @@
+{
+  "evaluations": [
+    {
+      "id": "sol-percent-of-peak",
+      "skill_should_trigger": true,
+      "prompt": "how close to peak is this kernel running in my rocprof-compute workload?\n\nThe application is ./vcopy.cpp.",
+      "workspace": "../rocprof-compute-sample",
+      "expected_behavior": [
+        "analyzes the sol block for a single kernel id",
+        "reads the Percent of Peak column from the report",
+        "compares against the Peak value the report prints rather than a remembered threshold"
+      ],
+      "unexpected_behavior": [
+        "calculates a theoretical peak instead of using the one in the report"
+      ],
+      "logs_contain": ["-b sol"]
+    },
+    {
+      "id": "sol-overall-efficiency",
+      "skill_should_trigger": true,
+      "prompt": "give me an overall efficiency summary of this rocprofiler-compute capture\n\nThe application is ./vcopy.cpp.",
+      "workspace": "../rocprof-compute-sample",
+      "expected_behavior": [
+        "selects one kernel with -k before reading Speed-of-Light",
+        "summarizes using the printed peak and percent of peak values"
+      ],
+      "unexpected_behavior": [
+        "reports a Speed-of-Light table averaged across unrelated kernels"
+      ],
+      "logs_contain": ["-b sol"]
+    },
+    {
+      "id": "sol-mfma-zero",
+      "skill_should_trigger": true,
+      "prompt": "my rocprof-compute report shows MFMA FLOPs at zero, is that a bug?",
+      "expected_behavior": [
+        "explains that zero is expected when the kernel issues no matrix instructions",
+        "suggests confirming the instruction mix before treating it as a problem"
+      ],
+      "unexpected_behavior": [
+        "treats zero MFMA as a defect without checking whether the kernel should use matrix instructions"
+      ]
+    },
+    {
+      "id": "sol-threshold-pushback",
+      "skill_should_trigger": true,
+      "prompt": "my kernel is at 62% of peak memory bandwidth, is that good enough?",
+      "expected_behavior": [
+        "declines to apply a fixed pass or fail cutoff",
+        "suggests comparing against neighbouring kernels or a known-good baseline"
+      ],
+      "unexpected_behavior": [
+        "states a universal threshold such as 70% as if the tool defined one"
+      ]
+    },
+    {
+      "id": "sol-instruction-level",
+      "skill_should_trigger": false,
+      "note": "Peak comparison is not the question; this is per-instruction work.",
+      "prompt": "which individual instruction in this kernel stalls the most?"
+    },
+    {
+      "id": "sol-cuda-peak",
+      "skill_should_trigger": false,
+      "prompt": "what percent of peak is my CUDA kernel hitting in Nsight Compute?"
+    },
+    {
+      "id": "sol-hardware-spec-lookup",
+      "skill_should_trigger": false,
+      "note": "Asks for a datasheet number, not a reading from a profile.",
+      "prompt": "what is the theoretical peak FP64 throughput of an MI300X?"
+    }
+  ]
+}
diff --git a/skills/rocprof-compute-speed-of-light/evals/machine.yml b/skills/rocprof-compute-speed-of-light/evals/machine.yml
new file mode 100644
index 00000000..4d7efd45
--- /dev/null
+++ b/skills/rocprof-compute-speed-of-light/evals/machine.yml
@@ -0,0 +1,2 @@
+os: [Linux]
+labels: [mi300x, gpu, rocm]
diff --git a/skills/rocprof-compute-speed-of-light/skill-card.md b/skills/rocprof-compute-speed-of-light/skill-card.md
new file mode 100644
index 00000000..eac7e0e5
--- /dev/null
+++ b/skills/rocprof-compute-speed-of-light/skill-card.md
@@ -0,0 +1,13 @@
+# Skill Card
+
+## Description
+
+Reads the rocprof-compute System Speed-of-Light panel to show how close an AMD GPU kernel runs to hardware peak.
+
+## Owner
+
+rocprofiler-compute, a project in [rocm-systems](https://github.com/ROCm/rocm-systems/)
+
+## License
+
+MIT
diff --git a/skills/rocprof-compute-torch-trace/.federated.json b/skills/rocprof-compute-torch-trace/.federated.json
new file mode 100644
index 00000000..c1becf20
--- /dev/null
+++ b/skills/rocprof-compute-torch-trace/.federated.json
@@ -0,0 +1,10 @@
+{
+  "source": "rocm-rocm-systems",
+  "repo": "ROCm/rocm-systems",
+  "ref": "develop",
+  "commit": "8cce7054f17a77272e1ba73ef78f4f77cf056956",
+  "path": "projects/rocprofiler-compute/skills/rocprof-compute-torch-trace",
+  "license": "MIT",
+  "content_hash": "7a8aa8e901c4e9124ea14de152ab6a4aa0d36814043c92401f9fb9f95ebc03a7",
+  "imported_at": "2026-10-28T01:19:39Z"
+}
diff --git a/skills/rocprof-compute-torch-trace/SKILL.md b/skills/rocprof-compute-torch-trace/SKILL.md
new file mode 100644
index 00000000..a813f2cd
--- /dev/null
+++ b/skills/rocprof-compute-torch-trace/SKILL.md
@@ -0,0 +1,127 @@
+---
+name: rocprof-compute-torch-trace
+description: Attributes AMD GPU kernel performance counters back to the PyTorch or Triton operators that launched them, using rocprof-compute torch and triton tracing. Use when the user asks which PyTorch layer or operator is slow, wants to map GPU kernels to model code, profiles a training or inference script, or asks about torch.compile and Inductor kernels. Not for CUDA tools, Windows, non-Python applications, or user-authored ROCTx ranges.
+---
+
+# Map GPU kernels to PyTorch operators
+
+Profiling a PyTorch model with counters alone gives kernel names, which say
+little about the model. Torch tracing annotates the run so each kernel can be
+traced back to the operator that launched it.
+
+Torch and Triton tracing are experimental. Every command needs
+`--experimental`, and option names can change between releases, so run
+`rocprof-compute profile --experimental --help` first.
+
+## 1. Check the requirements
+
+PyTorch version coverage, the native collector versus reduced
+`TorchDispatchMode` coverage, and the abort when two ROCm installs load in
+one process are in the Torch trace section of
+[profile mode](https://github.com/ROCm/rocm-systems/blob/8cce7054f17a77272e1ba73ef78f4f77cf056956/projects/rocprofiler-compute/docs/how-to/profile/mode.rst). Check the loader message
+before treating a missing backward operator as missing GPU work.
+
+The workload must be a Python command or Python script; the tracer cannot
+annotate a compiled binary. The collector is a generic prebuilt
+`torch_trace_collector.so` and needs no PyTorch or Python development packages
+to build. The workload still needs Python-compatible `roctx` bindings. The
+full requirement list is in the Torch trace section of
+[profile mode](https://github.com/ROCm/rocm-systems/blob/8cce7054f17a77272e1ba73ef78f4f77cf056956/projects/rocprofiler-compute/docs/how-to/profile/mode.rst).
+
+## 2. Collect
+
+```bash
+rocprof-compute profile \
+    --experimental \
+    --torch-trace \
+    --name <workload_name> \
+    -- python <script.py> <args>
+```
+
+For Triton kernels launched through Triton's Python runtime (`@triton.jit`):
+
+```bash
+rocprof-compute profile --experimental --triton-trace --name <name> -- python <script.py>
+```
+
+`torch.compile` / Inductor fused kernels use PyTorch's static launcher, so they
+do not get `|triton` ranges. Attribute them with `--torch-trace`.
+
+The two combine, and `--ml-api-trace` turns on every supported framework
+backend at once.
+
+Tracing adds overhead on top of the usual counter replay. Profile a few
+iterations of a steady-state model rather than a full training run, and skip
+warmup. `ROCPROFCOMPUTE_ROCTX_DEEP_TENSOR_WRAPS=0` reduces instrumentation
+overhead by dropping the device and layout wraps; see
+[profile mode](https://github.com/ROCm/rocm-systems/blob/8cce7054f17a77272e1ba73ef78f4f77cf056956/projects/rocprofiler-compute/docs/how-to/profile/mode.rst) for what that costs you.
+
+Tracing cannot be combined with attaching to a running process, and cannot be
+combined with a PC-sampling-only profile.
+
+When counters are being spread thin across many short operator kernels,
+`--iteration-multiplexing kernel` collects more per pass.
+
+Profile preserves each pass's raw marker and counter CSVs with the
+`ml_api_trace` prefix.
+
+## 3. List the captured operators
+
+List when you need the operator names from this run; filtering does not
+require a prior list. Do not pass `--list-*-operators` and `--*-operator`
+on the same command; listing ignores the filter. These flags require
+`--output-format stdout` (the default):
+
+```bash
+rocprof-compute analyze \
+    --path ./workloads/<name>/<gpu_model> \
+    --experimental --list-torch-operators
+```
+
+This prints the operators as a call tree sorted by GPU duration, with
+`file:line` on each node that recorded a source location, plus kernel
+launch statistics, so it also shows which operators are worth opening.
+
+The Triton equivalent is `--list-triton-operators`.
+
+## 4. Filter to an operator
+
+```bash
+rocprof-compute analyze \
+    --path ./workloads/<name>/<gpu_model> \
+    --experimental --torch-operator '*relu*'
+```
+
+Patterns are shell-style globs, not regexes, matched against a `/`-separated
+operator hierarchy. `*relu*` contains relu, `*/aten::addmm` is an addmm leaf
+on the reconstructed path, and `all` or `'*'` matches everything. Several
+patterns can be given, space or comma separated.
+Quote patterns so the shell does not expand them.
+
+`--triton-operator` works the same way. `*rmsnorm_kernel*` matches the
+rmsnorm nodes on a Triton tree.
+
+Worked examples are in the PyTorch operator analysis section of
+[the analyze CLI guide](https://github.com/ROCm/rocm-systems/blob/8cce7054f17a77272e1ba73ef78f4f77cf056956/projects/rocprofiler-compute/docs/how-to/analyze/cli.rst).
+
+## 5. Read the result
+
+Attribution tells you which operator owns the kernel time. It does not tell you
+why that kernel is slow. Once an operator is identified, analyze its kernels
+like any other: Speed-of-Light first, then the memory or roofline path.
+
+An operator that is slow because it launches thousands of tiny kernels is a
+different problem from one slow kernel, and the launch statistics in the
+operator list are what distinguish the two.
+
+## 6. ROCTx scope
+
+rocprof-compute uses ROCTx internally: torch tracing injects ROCTx regions to
+annotate operators, which is what makes attribution possible. That is the only
+place ROCTx is involved.
+
+rocprof-compute does not support ROCTx-range-based profiling control, such as
+starting, stopping, or resuming collection at a user-authored ROCTx range.
+User-defined ranges in the marker CSV are labeled `user`. If the user wants
+ROCTx start/stop control, say so plainly and offer kernel filtering with
+`-k` or dispatch filtering with `-d` instead.
diff --git a/skills/rocprof-compute-torch-trace/evals/evals.json b/skills/rocprof-compute-torch-trace/evals/evals.json
new file mode 100644
index 00000000..42ac53b8
--- /dev/null
+++ b/skills/rocprof-compute-torch-trace/evals/evals.json
@@ -0,0 +1,87 @@
+{
+  "evaluations": [
+    {
+      "id": "tt-which-operator-is-slow",
+      "skill_should_trigger": true,
+      "prompt": "which PyTorch operator is launching the slow GPU kernels in my model?\n\nThe application is ./simple_net.py.",
+      "workspace": "../rocprof-compute-sample",
+      "expected_behavior": [
+        "collects with --experimental --torch-trace",
+        "lists the captured operators to find the slow ones",
+        "checks the documented PyTorch version requirement"
+      ],
+      "unexpected_behavior": [
+        "guesses an operator pattern without listing the captured operators first"
+      ],
+      "logs_contain": ["--torch-trace", "--list-torch-operators"]
+    },
+    {
+      "id": "tt-map-kernels-to-model-code",
+      "skill_should_trigger": true,
+      "prompt": "profile my training script and tell me which layer the GPU time belongs to\n\nThe application is ./simple_net.py.",
+      "workspace": "../rocprof-compute-sample",
+      "expected_behavior": [
+        "profiles the Python script with torch tracing enabled",
+        "advises profiling a few steady-state iterations rather than a full run"
+      ],
+      "logs_contain": ["--torch-trace"]
+    },
+    {
+      "id": "tt-triton-inductor-kernels",
+      "skill_should_trigger": true,
+      "prompt": "my model uses torch.compile, how do I attribute the generated Triton kernels?",
+      "expected_behavior": [
+        "uses --torch-trace, because --triton-trace does not capture kernels that torch.compile generates"
+      ],
+      "unexpected_behavior": [
+        "relies on --triton-trace alone to capture torch.compile kernels"
+      ],
+      "logs_contain": ["--torch-trace"]
+    },
+    {
+      "id": "tt-filter-by-operator",
+      "skill_should_trigger": true,
+      "prompt": "show me the counters for just the conv2d operators in this torch-trace workload",
+      "expected_behavior": [
+        "filters with --torch-operator",
+        "treats the pattern as a shell-style glob rather than a regex"
+      ],
+      "logs_contain": ["--torch-operator"]
+    },
+    {
+      "id": "tt-custom-roctx-unsupported",
+      "skill_should_trigger": true,
+      "prompt": "I added my own ROCTx ranges to my HIP application, profile and report on those ranges",
+      "expected_behavior": [
+        "says plainly that rocprof-compute does not attribute results to user-authored ROCTx ranges",
+        "offers kernel filtering with -k or dispatch filtering with -d instead"
+      ],
+      "unexpected_behavior": [
+        "claims rocprof-compute can profile, filter, or start and stop collection on user-authored ROCTx ranges"
+      ]
+    },
+    {
+      "id": "tt-non-python-binary",
+      "skill_should_trigger": true,
+      "prompt": "attribute GPU kernels to functions in my compiled C++ HIP binary",
+      "expected_behavior": [
+        "says that operator attribution requires a Python workload and cannot annotate a compiled binary",
+        "offers kernel filtering with -k or dispatch filtering with -d instead"
+      ],
+      "unexpected_behavior": [
+        "claims torch or triton tracing can attribute kernels to functions in a compiled binary"
+      ]
+    },
+    {
+      "id": "tt-cuda-pytorch",
+      "skill_should_trigger": false,
+      "prompt": "profile my PyTorch model on an NVIDIA GPU with the torch profiler"
+    },
+    {
+      "id": "tt-model-accuracy",
+      "skill_should_trigger": false,
+      "note": "PyTorch question with no performance profiling in it.",
+      "prompt": "why is my PyTorch model's validation accuracy dropping?"
+    }
+  ]
+}
diff --git a/skills/rocprof-compute-torch-trace/evals/machine.yml b/skills/rocprof-compute-torch-trace/evals/machine.yml
new file mode 100644
index 00000000..4d7efd45
--- /dev/null
+++ b/skills/rocprof-compute-torch-trace/evals/machine.yml
@@ -0,0 +1,2 @@
+os: [Linux]
+labels: [mi300x, gpu, rocm]
diff --git a/skills/rocprof-compute-torch-trace/skill-card.md b/skills/rocprof-compute-torch-trace/skill-card.md
new file mode 100644
index 00000000..02bdddc3
--- /dev/null
+++ b/skills/rocprof-compute-torch-trace/skill-card.md
@@ -0,0 +1,13 @@
+# Skill Card
+
+## Description
+
+Attributes AMD GPU kernel performance back to the PyTorch and Triton operators that launched them with rocprof-compute.
+
+## Owner
+
+rocprofiler-compute, a project in [rocm-systems](https://github.com/ROCm/rocm-systems/)
+
+## License
+
+MIT