feat(emulation): debug emulated GPU kernels with real ROCgdb

emulation/mirage/ctl/src/lib.rs

@@ -603,6 +603,18 @@ pub struct RunArgs {
     /// daemon (the default).
     #[arg(long = "in-process")]
     in_process: bool,
+    /// Debug the workload interactively under ROCgdb: the command is
+    /// wrapped as `rocgdb --args <command...>`, with kernel breakpoints
+    /// made pending so `break <kernel>` works before the GPU code object
+    /// loads at launch. Requires a profile whose container ships ROCgdb.
+    #[arg(long)]
+    gdb: bool,
+    /// Extra ROCgdb command to run before the program, injected as
+    /// `-ex <CMD>` (repeatable, applied in order). Implies `--gdb`.
+    /// Useful for scripted/batch debugging, e.g.
+    /// `--gdb-ex 'break add_one' --gdb-ex run --gdb-ex continue`.
+    #[arg(long = "gdb-ex", value_name = "CMD")]
+    gdb_ex: Vec<String>,
     /// The command and its arguments.
     #[arg(trailing_var_arg = true, required = true, allow_hyphen_values = true)]
     argv: Vec<String>,
@@ -1765,6 +1777,32 @@ fn split_argv(argv: &[String]) -> (String, Vec<String>) {
     (cmd, it.collect())
 }
 
+/// Wrap a workload argv so it runs under ROCgdb for GPU kernel debugging
+/// (`mirage run --gdb`). `["./add_one", "5"]` becomes
+/// `rocgdb -ex "set breakpoint pending on" [-ex EX]... --args ./add_one 5`.
+///
+/// Kernel breakpoints are made pending because GPU code objects load lazily at
+/// dispatch: a `break <kernel>` set before launch would otherwise be rejected.
+/// Extra `-ex` commands (`--gdb-ex`) are injected in order, before `--args`, so
+/// a run can be fully scripted. When any `-ex` command is present the session is
+/// non-interactive, so `--batch` is added: ROCgdb runs the commands and exits
+/// (a bare `--gdb`, with no `-ex`, stays interactive and reads the terminal).
+fn gdb_wrap_argv(argv: &[String], gdb_ex: &[String]) -> Vec<String> {
+    let mut wrapped = vec!["rocgdb".to_string()];
+    if !gdb_ex.is_empty() {
+        wrapped.push("--batch".to_string());
+    }
+    wrapped.push("-ex".to_string());
+    wrapped.push("set breakpoint pending on".to_string());
+    for ex in gdb_ex {
+        wrapped.push("-ex".to_string());
+        wrapped.push(ex.clone());
+    }
+    wrapped.push("--args".to_string());
+    wrapped.extend_from_slice(argv);
+    wrapped
+}
+
 /// Parse repeated `KEY=VALUE` pairs from the CLI into the env map
 /// used by [`ExecArgs`]. Rejects entries without an `=` so a typo
 /// surfaces immediately instead of silently being dropped.
@@ -2013,7 +2051,14 @@ async fn run_cmd<C: MirageCtl + 'static>(ctl: Arc<C>, a: RunArgs) -> anyhow::Res
             (def.id, true)
         }
     };
-    let (cmd, args) = split_argv(&a.argv);
+    // `--gdb` / `--gdb-ex` wrap the workload under ROCgdb for interactive GPU
+    // kernel debugging. `--gdb-ex` implies `--gdb`.
+    let effective_argv = if a.gdb || !a.gdb_ex.is_empty() {
+        gdb_wrap_argv(&a.argv, &a.gdb_ex)
+    } else {
+        a.argv.clone()
+    };
+    let (cmd, args) = split_argv(&effective_argv);
     let env = parse_envs(&a.envs)?;
     let def = ExecDef {
         timestamp: chrono::Utc::now(),
@@ -2234,6 +2279,45 @@ mod tests {
         assert!(parse_option("=value").is_err());
     }
 
+    #[test]
+    fn gdb_wrap_prepends_rocgdb_with_pending_breakpoints() {
+        let got = gdb_wrap_argv(&["./add_one".to_string(), "5".to_string()], &[]);
+        assert_eq!(
+            got,
+            vec![
+                "rocgdb",
+                "-ex",
+                "set breakpoint pending on",
+                "--args",
+                "./add_one",
+                "5",
+            ]
+        );
+    }
+
+    #[test]
+    fn gdb_wrap_injects_extra_ex_commands_in_order() {
+        let got = gdb_wrap_argv(
+            &["./add_one".to_string()],
+            &["break add_one".to_string(), "run".to_string()],
+        );
+        assert_eq!(
+            got,
+            vec![
+                "rocgdb",
+                "--batch",
+                "-ex",
+                "set breakpoint pending on",
+                "-ex",
+                "break add_one",
+                "-ex",
+                "run",
+                "--args",
+                "./add_one",
+            ]
+        );
+    }
+
     #[test]
     fn no_overrides_keeps_by_name_ref() {
         let mut p = sample_profile();

emulation/mirage/rocjitsu_sys/src/lib.rs

@@ -115,6 +115,10 @@ pub struct RjVmCmd {
     /// In daemon mode the VM substitutes it into DBG_TRAP ENABLE and, on
     /// adoption, clears it to -1 so the caller does not close it.
     pub in_handle: RjHandle,
+    /// `[in/out]` Debugger-authorized target `/proc/pid/mem` fd, or -1.
+    pub in_mem_handle: RjHandle,
+    /// `[in]` Pinned target `/proc/pid` directory fd, or -1.
+    pub in_proc_handle: RjHandle,
 }
 
 /// Device memory mapping descriptor (`rj_vm_map_t`).
@@ -610,7 +614,7 @@ mod tests {
         assert_eq!(std::mem::size_of::<RjVmUnmap>(), 16);
         // rj_vm_cmd_t: u32 + (pad) + ptr + usize + i32 + i32 + i32 + (pad)
         // on 64-bit.
-        assert_eq!(std::mem::size_of::<RjVmCmd>(), 40);
+        assert_eq!(std::mem::size_of::<RjVmCmd>(), 48);
         assert_eq!(RjVmMode::Daemon as i32, 2);
         // rj_vm_gpu_info_t — must match the 312-byte RpcGpuInfo the
         // daemon handshake embeds (static_assert in rpc.h).

emulation/mirage/scripts/record_demo.sh

@@ -0,0 +1,116 @@
+#!/usr/bin/env bash
+#
+# Copyright (c) 2026 Advanced Micro Devices, Inc.
+# SPDX-License-Identifier: MIT
+#
+# Record an asciinema demo of a mirage/rocjitsu workflow.
+#
+# WHY: the emulation demos under `emulation/rocjitsu/demos/` each ship three
+# files — `<demo>.md` (what it shows), `<demo>.sh` (the exact commands), and
+# `<demo>.cast` (the recording). This helper is the single, reproducible way to
+# (re)generate the `.cast`: it builds a portable mirage + rocjitsu with
+# `mirage-docker-build.sh`, puts that `mirage` on `PATH`, and records the demo
+# script with asciinema, overwriting the `.cast` next to the `.sh` by default.
+#
+# Usage:
+#   scripts/record_demo.sh [options] <demo.sh> [output.cast]
+#
+# Examples:
+#   # Build mirage+rocjitsu, record emulation/rocjitsu/demos/rocgdb-quickstart.cast
+#   scripts/record_demo.sh ../rocjitsu/demos/rocgdb-quickstart.sh
+#
+#   # Reuse an already-built prefix (skip the ~minutes-long docker build)
+#   scripts/record_demo.sh --no-build --prefix ./build/manylinux \
+#       ../rocjitsu/demos/rocgdb-quickstart.sh
+#
+# Options:
+#   --no-build            Skip the docker build; reuse --prefix / $MIRAGE_BIN.
+#   --prefix DIR          Build into / read the mirage prefix here
+#                         (default: mirage/build/manylinux; also $MIRAGE_PREFIX).
+#   --title STR           asciinema recording title (default: the demo name).
+#   -h, --help            Show this help.
+#
+# Environment:
+#   MIRAGE_BIN            Explicit `mirage` binary to use (implies --no-build).
+#   MIRAGE_PREFIX         Same as --prefix.
+#   CARGO_PROFILE         Passed to mirage-docker-build.sh (default: release).
+#   ASCIINEMA_COLS/ROWS   Terminal size for the recording (default 100x30).
+
+set -euo pipefail
+
+SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
+MIRAGE_DIR="$(cd "$SCRIPT_DIR/.." && pwd)"
+
+no_build=0
+prefix="${MIRAGE_PREFIX:-$MIRAGE_DIR/build/manylinux}"
+title=""
+demo=""
+out=""
+
+usage() { sed -n '2,40p' "${BASH_SOURCE[0]}" | sed 's/^# \{0,1\}//'; }
+
+while [[ $# -gt 0 ]]; do
+  case "$1" in
+    --no-build) no_build=1; shift ;;
+    --prefix) prefix="$2"; shift 2 ;;
+    --title) title="$2"; shift 2 ;;
+    -h|--help) usage; exit 0 ;;
+    --) shift; break ;;
+    -*) echo "record_demo: unknown option: $1" >&2; usage >&2; exit 2 ;;
+    *) if [[ -z "$demo" ]]; then demo="$1"; elif [[ -z "$out" ]]; then out="$1"; else
+         echo "record_demo: unexpected argument: $1" >&2; exit 2; fi; shift ;;
+  esac
+done
+
+[[ -n "$demo" ]] || { echo "record_demo: missing <demo.sh>" >&2; usage >&2; exit 2; }
+[[ -f "$demo" ]] || { echo "record_demo: no such demo script: $demo" >&2; exit 1; }
+demo="$(cd "$(dirname "$demo")" && pwd)/$(basename "$demo")"
+# Default output: <demo>.cast next to the .sh, overwritten.
+[[ -n "$out" ]] || out="${demo%.sh}.cast"
+[[ -n "$title" ]] || title="$(basename "${demo%.sh}")"
+
+# --- Require asciinema ------------------------------------------------------
+# Recording a demo is a developer convenience, which is not a good enough reason
+# to mutate the machine it runs on. Provision the tool in the build or container
+# environment instead; this script only reports what is missing.
+command -v asciinema >/dev/null 2>&1 || {
+  echo "record_demo: asciinema not found. Install it and re-run, e.g." >&2
+  echo "  apt-get install asciinema  |  dnf install asciinema" >&2
+  echo "  pipx install asciinema     |  pip install --user asciinema" >&2
+  exit 1
+}
+
+# --- Build mirage + rocjitsu (portable) -------------------------------------
+mirage_bin="${MIRAGE_BIN:-}"
+if [[ -z "$mirage_bin" && "$no_build" -eq 0 ]]; then
+  echo "record_demo: building mirage + rocjitsu into $prefix" >&2
+  "$SCRIPT_DIR/mirage-docker-build.sh" "$prefix" >&2
+fi
+if [[ -z "$mirage_bin" ]]; then
+  mirage_bin="$prefix/bin/mirage"
+fi
+[[ -x "$mirage_bin" ]] || {
+  echo "record_demo: mirage binary not found at $mirage_bin" >&2
+  echo "  build it first (drop --no-build) or set \$MIRAGE_BIN." >&2
+  exit 1
+}
+mirage_bin="$(cd "$(dirname "$mirage_bin")" && pwd)/$(basename "$mirage_bin")"
+echo "record_demo: using mirage: $mirage_bin" >&2
+
+# --- Record -----------------------------------------------------------------
+# The demo script finds `mirage` on PATH (built above) and may read $MIRAGE_BIN.
+# asciinema replays the exact stdout of `bash <demo.sh>`; --overwrite keeps the
+# `.cast` regenerable in place.
+echo "record_demo: recording $demo -> $out" >&2
+env \
+  PATH="$(dirname "$mirage_bin"):$PATH" \
+  MIRAGE_BIN="$mirage_bin" \
+  asciinema rec \
+    --overwrite \
+    --title "$title" \
+    --cols "${ASCIINEMA_COLS:-100}" \
+    --rows "${ASCIINEMA_ROWS:-30}" \
+    --command "bash '$demo'" \
+    "$out"
+
+echo "record_demo: wrote $out" >&2

emulation/rocjitsu/README.md

@@ -122,6 +122,7 @@ See [docs/building.md](docs/building.md) for container setup with PyTorch.
 | [Configuration](docs/configuration.md) | JSON config format and topology |
 | [CLI & Transport](docs/rocjitsu-cli.md) | CLI modes, daemon RPC protocol |
 | [Race Detector](docs/race-detector.md) | Race detection tutorial and internals |
+| [Debugging with ROCgdb](docs/rocgdb-debugging.md) | Debug emulated GPU kernels with ROCgdb: breakpoints, watchpoints, faults, multi-wave |
 
 ### Design
 

emulation/rocjitsu/demos/rocgdb-fault.cast

@@ -0,0 +1,17 @@
+{"version": 2, "width": 100, "height": 30, "timestamp": 1783636162, "env": {"SHELL": "/bin/bash", "TERM": "xterm-256color"}, "title": "rocgdb-fault"}
+[0.009987, "o", "\r\n\u001b[1;36m# A HIP kernel that stores through a wild, never-mapped device pointer:\u001b[0m\r\n"]
+[0.011269, "o", "__global__ void bad_access(int *data, int n) {\r\n  int i = blockIdx.x * blockDim.x + threadIdx.x;\r\n  // A wild, never-mapped device address. The store below faults.\r\n  int *wild = reinterpret_cast<int *>(0x0000dead0000ULL);\r\n  if (i < n)\r\n    wild[i] = data[i] + 1;\r\n}\r\n"]
+[0.011443, "o", "\r\n\u001b[1;36m# Build it for gfx950 (MI350X) with device debug info:\u001b[0m\r\n"]
+[0.014641, "o", "+ hipcc --offload-arch=gfx950 -g -O0 -o /tmp/tmp.wtWfqkbt4t/bad_access /home/arosa/rocm-systems/emulation/rocjitsu/demos/../tests/rocgdb/bad_access.hip\r\n"]
+[0.531921, "o", "\u001b[1m/home/arosa/rocm-systems/emulation/rocjitsu/demos/../tests/rocgdb/bad_access.hip:29:3: \u001b[0m\u001b[0;1;35mwarning: \u001b[0m\u001b[1mignoring return value of type 'hipError_t' declared with 'nodiscard' attribute [-Wunused-value]\u001b[0m\r\n   29 |   hipFree(d);\u001b[0m\r\n      | \u001b[0;1;32m  ^~~~~~~~~~\r\n\u001b[0m"]
+[0.635648, "o", "1 warning generated when compiling for gfx950.\r\n"]
+[1.136021, "o", "\u001b[1m/home/arosa/rocm-systems/emulation/rocjitsu/demos/../tests/rocgdb/bad_access.hip:29:3: \u001b[0m\u001b[0;1;35mwarning: \u001b[0m\u001b[1mignoring return value of type 'hipError_t' declared with 'nodiscard' attribute [-Wunused-value]\u001b[0m\r\n   29 |   hipFree(d);\u001b[0m\r\n      | \u001b[0;1;32m  ^~~~~~~~~~\r\n\u001b[0m"]
+[1.189941, "o", "1 warning generated when compiling for host.\r\n"]
+[1.221753, "o", "\r\n\u001b[1;36m# Run it under ROCgdb — the emulator faults the wave on the wild store:\u001b[0m\r\n"]
+[1.222347, "o", "+ /home/arosa/rocm-systems/emulation/mirage/target/debug/mirage run --profile mi350x --gdb --gdb-ex run --gdb-ex backtrace --gdb-ex 'info registers pc' -- /tmp/tmp.wtWfqkbt4t/bad_access\r\n"]
+[2.484295, "o", "[Thread debugging using libthread_db enabled]\r\nUsing host libthread_db library \"/lib/x86_64-linux-gnu/libthread_db.so.1\".\r\n"]
+[2.633349, "o", "[New Thread 0x7ffff27ff6c0 (LWP 1839316)]\r\n[New Thread 0x7fffebfff6c0 (LWP 1839317)]\r\n[New Thread 0x7ffff1ffe6c0 (LWP 1839319)]\r\n[Thread 0x7ffff1ffe6c0 (LWP 1839319) exited]\r\n"]
+[2.883424, "o", "[New Thread 0x7ffff17fd6c0 (LWP 1839399)]\r\n[Thread 0x7ffff17fd6c0 (LWP 1839399) exited]\r\n[New Thread 0x7ffff2d7e6c0 (LWP 1839400)]\r\n"]
+[2.983503, "o", "\r\nThread 7 received signal SIGSEGV, Segmentation fault.\r\nWarning: precise memory violation signal reporting is not enabled, reported\r\nlocation may not be accurate.  See \"show amdgpu precise-memory\".\r\n[Switching to thread 7, lane 0 (AMDGPU Lane 1:2:?:1/0 (?,?,?)[?,?,?])]\r\n"]
+[3.034011, "o", "0x000000540002d11c in bad_access (data=<error reading variable: Cannot access memory at address private_lane#0x20>, n=<error reading variable: Cannot access memory at address private_lane#0x28>) at /home/arosa/rocm-systems/emulation/rocjitsu/demos/../tests/rocgdb/bad_access.hip:17\r\n17\t    wild[i] = data[i] + 1;\r\n#0  0x000000540002d11c in bad_access (data=<error reading variable: Cannot access memory at address private_lane#0x20>, n=<error reading variable: Cannot access memory at address private_lane#0x28>) at /home/arosa/rocm-systems/emulation/rocjitsu/demos/../tests/rocgdb/bad_access.hip:17\r\npc             0x540002d11c        0x540002d11c <bad_access(int*, int)+1052>\r\n"]
+[4.097695, "o", "\r\n\u001b[1;36m# ROCgdb caught the GPU memory violation as SIGSEGV, stopped at the faulting store — no AMD GPU required.\u001b[0m\r\n"]

emulation/rocjitsu/demos/rocgdb-fault.md

@@ -0,0 +1,31 @@
+# rocgdb-fault
+
+Catch a GPU **memory-access fault (SIGSEGV)** with ROCgdb — in software, with no
+AMD GPU required.
+
+A wave that dereferences an unmapped device address is the GPU equivalent of a
+segfault. On real hardware it raises a memory-access fault that stops the wave;
+on the emulated MI350X the memory pipeline detects the unmapped access, raises
+`TRAPSTS.xnack_error`, and reports it to rocm-dbgapi as
+`WAVE_STOP_REASON_MEMORY_VIOLATION`, which ROCgdb surfaces as **SIGSEGV** at the
+faulting instruction.
+
+This demo builds a tiny HIP kernel that stores through a wild pointer
+(`0x0000dead0000`) and runs it under `mirage run --gdb`: ROCgdb stops the GPU
+wave with a segmentation fault at the offending store, so you can see the PC and
+backtrace of the wave that went out of bounds.
+
+## Run it
+
+Record (builds a portable mirage + rocjitsu, then captures the `.cast`):
+
+    emulation/mirage/scripts/record_demo.sh emulation/rocjitsu/demos/rocgdb-fault.sh
+
+Or run against an already-built `mirage`:
+
+    MIRAGE_BIN=/path/to/mirage bash emulation/rocjitsu/demos/rocgdb-fault.sh
+
+Requires `hipcc` and `rocgdb` on `PATH` (from a ROCm install).
+
+See [rocgdb-debugging.md](../docs/rocgdb-debugging.md) for how the emulator maps
+GPU exceptions to debugger stop reasons.

emulation/rocjitsu/demos/rocgdb-fault.sh

@@ -0,0 +1,36 @@
+#!/usr/bin/env bash
+#
+# Copyright (c) 2026 Advanced Micro Devices, Inc.
+# SPDX-License-Identifier: MIT
+#
+# Demo: catch a GPU memory-access fault (SIGSEGV) with ROCgdb, entirely in
+# software with no AMD GPU. See rocgdb-fault.md. Regenerate the .cast with:
+#   emulation/mirage/scripts/record_demo.sh emulation/rocjitsu/demos/rocgdb-fault.sh
+
+set -euo pipefail
+
+here="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
+kernel="$here/../tests/rocgdb/bad_access.hip"
+mirage="${MIRAGE_BIN:-mirage}"
+
+say() { printf '\n\033[1;36m# %s\033[0m\n' "$*"; }
+
+say "A HIP kernel that stores through a wild, never-mapped device pointer:"
+sed -n '/__global__/,/^}/p' "$kernel"
+
+say "Build it for gfx950 (MI350X) with device debug info:"
+tmp="$(mktemp -d)"
+trap 'rm -rf "$tmp"' EXIT
+app="$tmp/bad_access"
+( set -x; hipcc --offload-arch=gfx950 -g -O0 -o "$app" "$kernel" )
+
+say "Run it under ROCgdb — the emulator faults the wave on the wild store:"
+( set -x
+  "$mirage" run --profile mi350x --gdb \
+    --gdb-ex 'run' \
+    --gdb-ex 'backtrace' \
+    --gdb-ex 'info registers pc' \
+    -- "$app"
+)
+
+say "ROCgdb caught the GPU memory violation as SIGSEGV, stopped at the faulting store — no AMD GPU required."