diff --git a/.github/ISSUE_TEMPLATE/bug_report.md b/.github/ISSUE_TEMPLATE/bug_report.md new file mode 100644 index 0000000..96aa435 --- /dev/null +++ b/.github/ISSUE_TEMPLATE/bug_report.md @@ -0,0 +1,11 @@ +--- +name: Bug report +about: Something broke on your AMD GPU +title: "bug: " +labels: bug +--- +**Hardware / stack** (paste `scripts/00-check-env.sh` output or +docs/results/environment.json): +**Exact command:** +**Full log** (pastebin/gist if long): +**What you expected vs what happened:** diff --git a/.github/ISSUE_TEMPLATE/validation_report.md b/.github/ISSUE_TEMPLATE/validation_report.md new file mode 100644 index 0000000..6da2317 --- /dev/null +++ b/.github/ISSUE_TEMPLATE/validation_report.md @@ -0,0 +1,11 @@ +--- +name: Validation report +about: Share results from YOUR hardware (successes AND failures) +title: "validation: / ROCm " +labels: validation +--- +**GPU / ROCm / torch:** +**`00-check-env.sh` output:** +**Summary table** (`python3 scripts/summarize_results.py`): +**Receipts** (docs/results/… if you can share): +**Anything that failed, with the log:** diff --git a/.github/PULL_REQUEST_TEMPLATE.md b/.github/PULL_REQUEST_TEMPLATE.md new file mode 100644 index 0000000..d6527c8 --- /dev/null +++ b/.github/PULL_REQUEST_TEMPLATE.md @@ -0,0 +1,10 @@ +**What & why** (one paragraph; link the spec/design doc if any): + +**Evidence** (receipts under docs/results/, or `verified on: ` + +link for ports from sibling repos — see CONTRIBUTING): + +**Gates** (must both be green): +- [ ] `python3 -m pytest tests/ -q` +- [ ] `shellcheck scripts/*.sh scripts/lib/common.sh docs/hardware/strix-halo/hsa_fix.sh` + +**If ported from SenseNova-U1.5-ROCm-8060S:** source commit SHA: diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 4a2bec6..8105209 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -19,4 +19,4 @@ jobs: - name: ShellCheck the pipeline scripts run: | sudo apt-get install -y shellcheck -q - shellcheck scripts/*.sh scripts/lib/common.sh + shellcheck scripts/*.sh scripts/lib/common.sh docs/hardware/strix-halo/hsa_fix.sh diff --git a/CHANGELOG.md b/CHANGELOG.md index 1781c23..a694ab8 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -16,6 +16,28 @@ and the project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0. to PR head e2f2c865; V1–V4 on-model verification (list input fixed, tuple output byte-identical) — receipts docs/results/validation/interleave-image-size/. +- Patch 0003: optional torch.compile/cudagraph safety (env-gated + `SENSENOVA_COMPILE*`); ported from the Strix Halo sibling repo — + verified on gfx1151 (V1–V8 + quality bench there); gfx1100 receipt + pending, compile stays opt-in. +- `scripts/gpu_monitor.py`: sample VRAM + GTT while a workload runs + (JSONL + peak summary) — GTT visibility is what evidences APU + GTT-spill runs. +- `scripts/rocm_check.py`: env sanity check (HIP build, wheel gfx + coverage, bf16 matmul + SDPA smoke, optional GTT-spill alloc test). +- Hardware profile `docs/hardware/strix-halo/` (EN+CN): GTT-spill + memory model, HSA runtime shim (`hsa_fix.sh`, now CI-shellchecked), + gfx1151 wheel guidance, AOTriton hd72 caveat — the umbrella's first + APU profile, ported from the 8060S repo. +- `docs/porting.md` + `docs/porting_CN.md`: generic ROCm porting notes, + split from the 8060S repo's porting doc (APU-specific depth moved to + the strix-halo profile). +- `docs/upstream/aotriton-54.md`: upstream tracking entry for the + AOTriton hd=72 silent-wrong-output bug (issue #54, fix unmerged); + verified on gfx1151, evidence in the 8060S repo. +- Governance: Code of Conduct, `SECURITY.md`, bug/validation issue + templates, PR template (evidence gates), CONTRIBUTING cross-repo + section (graduation model), README/README_CN hardware-coverage table. ### Changed - pytorch#194498 root cause verified on-host: SDPA fused-backend launch diff --git a/CODE_OF_CONDUCT.md b/CODE_OF_CONDUCT.md new file mode 100644 index 0000000..123b278 --- /dev/null +++ b/CODE_OF_CONDUCT.md @@ -0,0 +1,22 @@ +# Code of Conduct + +## Our pledge +Be evidence-first about claims and kind to people. We welcome +contributors of every background and skill level. + +## Standards +**Encouraged:** citing receipts for claims, asking questions, reporting +both successes and failures, reviewing others' work respectfully. +**Unacceptable:** harassment or discriminatory language, hostile +responses to newcomers, presenting unverified numbers as measured, +doxxing, spam. + +## Enforcement +Report violations by opening a private GitHub security advisory +("Report a vulnerability" → select "not applicable") or contacting a +maintainer via GitHub. Maintainers will respond within 7 days and may +remove comments/commits/ban repeat offenders. + +## Attribution +Adapted in spirit from the Contributor Covenant 2.1. +本行为准则同样适用于中文社区讨论；违反行为可同样通过上述渠道举报。 diff --git a/CONTRIBUTING.md b/CONTRIBUTING.md index 217668d..1f5ff84 100644 --- a/CONTRIBUTING.md +++ b/CONTRIBUTING.md @@ -51,3 +51,25 @@ re-runs a single block for a minimal reproduction. - Scripts: `set -euo pipefail`, a `usage()` with `--help`, clear `[project]` log prefixes — follow the existing scripts. - Numbers in docs: cite the receipt file they came from. + +## Contributing between the two repos + +This repo is the **umbrella** (hardware-generic assets + hardware +profiles). [SenseNova-U1.5-ROCm-8060S](https://github.com/AIwork4me/SenseNova-U1.5-ROCm-8060S) is the **Strix Halo lab**: +APU-first experiments live there; verified work graduates here. + +**8060S → umbrella (graduation).** A change graduates when it is +① hardware-independent or absorbed as a hardware profile, ② verified on +gfx1151 with evidence in the 8060S repo, ③ passes this repo's offline +CI, ④ documented EN+CN, ⑤ not a duplicate of an existing tool here. +Open a PR labelled `backport`; the commit carries +`Ported from …@` and the claim is labelled +`verified on: gfx1151` until a gfx1100 receipt lands here. + +**Umbrella → 8060S (flow-down).** Changes to shared surfaces +(`patches/`, `scripts/`) are cherry-picked by the 8060S repo as +needed; that repo records `Merged-upstream: `. + +**Evidence across repos.** `receipt.py` records `hardware=` on +every run; cross-hardware claims must link the producing repo's +evidence. diff --git a/README.md b/README.md index d9a7504..ab69e01 100644 --- a/README.md +++ b/README.md @@ -36,6 +36,17 @@ bash scripts/quickstart.sh # t2i | vqa | edit | interleave [Troubleshooting](docs/getting-started.md#troubleshooting) · [Contributing](CONTRIBUTING.md) +## Hardware coverage + +| hardware | status | where | +|---|---|---| +| gfx1100 (RDNA3 dGPU) | fully validated — receipts in [docs/results/](docs/results/) | this repo | +| gfx1151 (Strix Halo APU) | fully validated — evidence in the [8060S repo](https://github.com/AIwork4me/SenseNova-U1.5-ROCm-8060S); profile: [docs/hardware/strix-halo/](docs/hardware/strix-halo/) | this repo + Strix Halo lab | + +Other AMD GPUs: the pipeline is hardware-generic (see +[Porting Notes](docs/porting.md)) — validation reports welcome +(issue template `validation_report`). + ## Showcase: movie-poster style Generated on the reference gfx1100 host with this repo's pipeline diff --git a/README_CN.md b/README_CN.md index fdc8be6..897c0fc 100644 --- a/README_CN.md +++ b/README_CN.md @@ -27,11 +27,25 @@ [SenseNova-U1#260](https://github.com/OpenSenseNova/SenseNova-U1/pull/260)）与 ROCm SDPA 融合后端崩溃（补丁 0002，[pytorch#194498](https://github.com/pytorch/pytorch/issues/194498)、 [SenseNova-U1#261](https://github.com/OpenSenseNova/SenseNova-U1/issues/261)） +- 🧪 **gfx1151（Strix Halo APU）同样全量验证**：证据见姊妹仓库 + [SenseNova-U1.5-ROCm-8060S](https://github.com/AIwork4me/SenseNova-U1.5-ROCm-8060S) + （Strix Halo 实验室），硬件配置在本仓库 + [docs/hardware/strix-halo/](docs/hardware/strix-halo/) ```bash bash scripts/quickstart.sh # 可选任务: t2i | vqa | edit | interleave ``` +## 硬件覆盖 + +| 硬件 | 状态 | 位置 | +|---|---|---| +| gfx1100（RDNA3 独显） | 完全验证 —— 收据见 [docs/results/](docs/results/) | 本仓库 | +| gfx1151（Strix Halo APU） | 完全验证 —— 证据见 [8060S 仓库](https://github.com/AIwork4me/SenseNova-U1.5-ROCm-8060S)；配置：[docs/hardware/strix-halo/](docs/hardware/strix-halo/) | 本仓库 + Strix Halo 实验室 | + +其他 AMD GPU：本管线与硬件无关（见[移植指南](docs/porting.md)）—— +欢迎提交验证报告（issue 模板 `validation_report`）。 + ## 效果展示：电影海报风格 以下海报由本仓库管线在 gfx1100 参考机上生成 diff --git a/SECURITY.md b/SECURITY.md new file mode 100644 index 0000000..75e9eb6 --- /dev/null +++ b/SECURITY.md @@ -0,0 +1,17 @@ +# Security Policy + +## Supported versions +The tip of `main` only. + +## Reporting a vulnerability +Please do NOT open a public issue. Use GitHub's private vulnerability +reporting ("Security" tab → "Report a vulnerability"). Include +reproduction steps and affected components (scripts/, patches/, docs/ +pipeline). Expect a response within 7 days. + +## Scope +This repo ships shell/python scripts and upstream patches; it does not +serve anything to the network. Model weights come from upstream +ModelScope — weight integrity is verified by +configs/artifact-manifest.json SHA256s (any tampering report is +welcome but belongs upstream). diff --git a/docs/getting-started.md b/docs/getting-started.md index a146fc7..5a3db5e 100644 --- a/docs/getting-started.md +++ b/docs/getting-started.md @@ -2,6 +2,9 @@ From a bare Linux host with an AMD GPU to your first SenseNova-U1.5 output. +For how this port works (wheels, attention backend, memory tiers): +[Porting Notes](porting.md) · [移植笔记](porting_CN.md) + ## What you need | Requirement | Why | Reference host | diff --git a/docs/hardware/strix-halo/README.md b/docs/hardware/strix-halo/README.md new file mode 100644 index 0000000..00b2369 --- /dev/null +++ b/docs/hardware/strix-halo/README.md @@ -0,0 +1,93 @@ +# Hardware Profile: Strix Halo (Radeon 8060S iGPU, gfx1151) + +verified on: gfx1151 — Ryzen AI MAX+ PRO 395, ROCm 7.2.1 system + +torch 2.12.0+rocm7.14.0 wheels. Full evidence corpus (V1–V8 speedup, +quality bench, determinism, tier tests): +[SenseNova-U1.5-ROCm-8060S](https://github.com/AIwork4me/SenseNova-U1.5-ROCm-8060S). +gfx1100 receipts live in [docs/results/](../../results/) here. +Generic background: [porting notes](../../porting.md). + +## Why full bf16 fits a "32 GB" iGPU — GTT spill + +On AMD APUs the GPU shares physical memory with the CPU. The driver +exposes two pools: dedicated **VRAM** (BIOS carve-out, e.g. 32 GB on the +reference machine) and **GTT** (the rest of system memory the GPU may +map, e.g. 80 GB). ROCm's allocator can allocate past dedicated VRAM by +spilling into GTT, so `torch.empty(50 GiB, device="cuda")` succeeds. +Since both pools are the same physical LPDDR5X on package, the penalty +is modest — in the 8060S repo's suite runs the allocator put ~1.2 GiB in +VRAM and up to 41.6 GiB in GTT (112 GiB total GPU-addressable), with +full-mode peaks ≤ 42.9 GiB, so ~45+ GiB addressable works in practice. + +`scripts/rocm_check.py --alloc-test 48` verifies exactly this spill +(48 GiB allocation; archived as `evidence/stack714/env_check.json` in +the 8060S repo). This is what makes `--vram_mode full` (entire ~50 GB +checkpoint GPU-resident, no layer offload) possible on a "32 GB" GPU. +Note when choosing tiers on the 7.14 stack: `full` is also the fastest +tier there — do not override to `fast`/`balanced` for speed (8060S +PERFORMANCE.md, 2026-08-25 correction; tiers remain the tool for VRAM +capacity on smaller GPUs). + +## HSA runtime pitfall (torch +rocm7.0 wheels, driver ≥ 7.1) + +Official `torch 2.10.0+rocm7.0` wheels bundle a ROCm 7.0 HSA runtime +that segfaults on gfx1151 (Strix Halo) when the kernel driver is from +ROCm ≥ 7.1: crash in `GpuAgent::QueueCreate` → `ReleaseQueueMainScratch` +while freezing the first code object — every GPU op dies. Fix: preload +the system ROCm runtime +(`LD_PRELOAD=/opt/rocm/lib/libhsa-runtime64.so`), which always matches +the kernel driver. + +This profile ships the shim as [hsa_fix.sh](hsa_fix.sh) — source it +(not execute) and call `hsa_fix_apply "$PY" "$REPO_ROOT"`; it exports +`LD_PRELOAD` only when the system runtime is NEWER than the wheel's +bundled one, so machines where the bundled runtime already works (or is +newer) are untouched. The recommended 7.14 wheels run bare and never +trigger it. + +## Wheel guidance for gfx1151 + +- **Default (recommended):** AMD official multi-arch + `torch 2.12.0+rocm7.14.0` from + `https://repo.amd.com/rocm/whl-multi-arch/` — verified bare-run on + gfx1151 (no HSA preload needed) and faster than the torch-2.10 build + in the 8060S A/B (167.9 s vs 232.1 s for a 1024×1024 t2i in `fast` + mode; both sessions were contended, so treat the magnitude as + indicative). These wheels bundle their own ROCm userspace — a system + ROCm mainly provides the kernel-driver tooling (`rocminfo`, + `rocm-smi`). +- **Fallback:** official PyTorch rocm7.0 index + (`https://download.pytorch.org/whl/rocm7.0`, `torch 2.10.0+rocm7.0`; + covers gfx942/950/1100/1101/1151/1200) — on gfx1151 it needs the HSA + preload above. +- **Last resort:** AMD arch-specific nightlies + (`https://rocm.nightlies.amd.com/v2//`) for arches the official + wheels miss. + +In this repo `scripts/01-setup-venv.sh` installs from a single index — +point `PY_INDEX_ROCM` at the source you want; the deeper +wheel-has-kernels-for-your-gfx check is `scripts/rocm_check.py` +(see [porting notes §4](../../porting.md#4-pytorch-wheel-matrix-known-rocm-wheel-sources)). + +## AOTriton caveat (hd=72 ViT heads) + +On gfx1151 with torch 2.12+rocm7.14, AOTriton flash SDPA is *silently +wrong* for head_dim=72 with seq>1024, **non-causal only** (causal +attention at the same shape is correct) — errors drift run-to-run, +typically 16–730% relative error, dominantly finite garbage with +sporadic NaN in ~1/5 of calls. SenseNova's own attention shapes +(head_dim 64/128) are unaffected, but disable AOTriton for a foreign +model with a 16-head/1152-hidden ViT. Root cause: upstream AOTriton +[issue #54](https://github.com/ROCm/aotriton/issues/54) (missing +`out_dtype=tl.float32` in the fwd kernel); a fix exists (commit +`8232d69672`) but is unmerged in every release so far — tracked in +[docs/upstream/aotriton-54.md](../../upstream/aotriton-54.md). + +## Companion repo + +The daily-driver run scripts, one-command quickstart, and the full +evidence corpus live in the +[8060S repo](https://github.com/AIwork4me/SenseNova-U1.5-ROCm-8060S) — +the Strix Halo lab where these numbers were measured. Experiments there +graduate into this umbrella via PR when they meet the graduation +criteria ([CONTRIBUTING.md](../../../CONTRIBUTING.md)). diff --git a/docs/hardware/strix-halo/README_CN.md b/docs/hardware/strix-halo/README_CN.md new file mode 100644 index 0000000..49320e8 --- /dev/null +++ b/docs/hardware/strix-halo/README_CN.md @@ -0,0 +1,82 @@ +# 硬件 Profile：Strix Halo（Radeon 8060S 核显，gfx1151） + +verified on: gfx1151 —— Ryzen AI MAX+ PRO 395，系统 ROCm 7.2.1 + +torch 2.12.0+rocm7.14.0 wheel。完整证据集（V1–V8 提速、质量评测、 +确定性、档位测试）： +[SenseNova-U1.5-ROCm-8060S](https://github.com/AIwork4me/SenseNova-U1.5-ROCm-8060S)。 +gfx1100 凭据在本仓库 [docs/results/](../../results/)。 +通用背景见[移植笔记](../../porting_CN.md)。 + +## 为什么"32 GB"核显装得下完整 bf16 —— GTT 溢出 + +AMD APU 的 GPU 与 CPU 共享物理内存。驱动暴露两个池：专用 **VRAM** +（BIOS 挖出，参考机型为 32 GB）与 **GTT**（GPU 可映射的其余系统内存， +如 80 GB）。ROCm 分配器越过专用 VRAM 时会溢出到 GTT，因此 +`torch.empty(50 GiB, device="cuda")` 能成功。两个池都是同一颗封装内 +的 LPDDR5X 物理内存，代价不大——8060S 仓库套件运行期间，分配器把 +~1.2 GiB 放进 VRAM、最多 41.6 GiB 进 GTT（GPU 可寻址合计 112 GiB）， +full 档峰值 ≤ 42.9 GiB，实践上 ~45+ GiB 可寻址即可。 + +本仓库的 `scripts/rocm_check.py --alloc-test 48` 正是验证这一溢出 +（48 GiB 分配；8060S 仓库存证 `evidence/stack714/env_check.json`）。 +这就是"32 GB 显存"也能跑 `--vram_mode full`（约 50 GB 权重全驻留 +GPU、不做 layer offload）的原因。7.14 栈上选档提示：`full` 同时也是 +最快档——不要为提速改选 `fast`/`balanced`（8060S PERFORMANCE.md +2026-08-25 更正；档位仍是小显存 GPU 的容量工具）。 + +## HSA 运行时的坑（torch +rocm7.0 wheel，驱动 ≥ 7.1） + +官方 `torch 2.10.0+rocm7.0` 轮子自带 ROCm 7.0 的 HSA 运行时，在 +gfx1151（Strix Halo）+ ROCm ≥ 7.1 内核驱动上冻结首个 code object 时 +（`GpuAgent::QueueCreate` → `ReleaseQueueMainScratch`）段错误——任何 +GPU 算子都跑不起来。解法：预加载系统 ROCm 运行时 +（`LD_PRELOAD=/opt/rocm/lib/libhsa-runtime64.so`），它必然与内核驱动 +匹配。 + +本 profile 以 [hsa_fix.sh](hsa_fix.sh) 携带该 shim——source（而非 +执行）本文件后调用 `hsa_fix_apply "$PY" "$REPO_ROOT"`；仅当系统运行时 +比 wheel 自带的**更新**时才 export `LD_PRELOAD`，自带运行时本来就 +正常（或更新）的机器不受影响。推荐的 7.14 wheel 裸跑，不会触发。 + +## gfx1151 wheel 指南 + +- **默认（推荐）**：AMD 官方多架构 `torch 2.12.0+rocm7.14.0`， + 索引 `https://repo.amd.com/rocm/whl-multi-arch/`——gfx1151 实测 + 裸跑（无需 HSA 预加载），且在 8060S A/B 中快于 torch-2.10 构建 + （1024×1024 t2i fast 档 167.9 s vs 232.1 s——当年会话存在争用， + 幅度仅供参考）。这类 wheel 自带 ROCm 用户态——系统 ROCm 主要提供 + 内核驱动工具（`rocminfo`、`rocm-smi`）。 +- **回退**：官方 PyTorch rocm7.0 索引 + （`https://download.pytorch.org/whl/rocm7.0`，`torch 2.10.0+rocm7.0`； + 覆盖 gfx942/950/1100/1101/1151/1200）——gfx1151 上需上面的 HSA + 预加载。 +- **兜底**：AMD 按架构 nightly + （`https://rocm.nightlies.amd.com/v2//`），用于官方 wheel 缺 + 架构时。 + +本仓库的 `scripts/01-setup-venv.sh` 从单一索引安装——用 +`PY_INDEX_ROCM` 指向你想要的来源；更深入的"wheel 是否带你 gfx 内核" +检查是 `scripts/rocm_check.py` +（见[移植笔记 §4](../../porting_CN.md#4-pytorch-wheel-矩阵已知-rocm-wheel-来源)）。 + +## AOTriton 告警（hd=72 ViT 头） + +gfx1151 + torch 2.12+rocm7.14 下，AOTriton flash SDPA 在 head_dim=72 +且 seq>1024、**仅非因果注意力**时**静默输出错误值**（同形状因果注意力 +正常）——误差逐次漂移，典型相对误差 16~730%，以有限错误值为主、约 +1/5 次调用出现 NaN。SenseNova 自身的注意力形状（head_dim 64/128） +不受影响；但若复用环境跑 16 头/hidden 1152 的外部 ViT 模型，请对其 +关闭 AOTriton。根因：上游 AOTriton +[issue #54](https://github.com/ROCm/aotriton/issues/54)（fwd 内核缺 +`out_dtype=tl.float32`）；修复 commit `8232d69672` 已存在但至今未合并 +进任何发布——追踪详情见 +[docs/upstream/aotriton-54.md](../../upstream/aotriton-54.md)。 + +## 配套仓库 + +日常驱动脚本、一条命令 quickstart 与完整证据集都在 +[8060S 仓库](https://github.com/AIwork4me/SenseNova-U1.5-ROCm-8060S) +——即 Strix Halo 实验室，上述数字都在那里测得。实验在那里达到毕业 +标准（graduation criteria， +[CONTRIBUTING.md](../../../CONTRIBUTING.md)）后，通过 PR 毕业进入 +本 umbrella 仓库。 diff --git a/docs/hardware/strix-halo/hsa_fix.sh b/docs/hardware/strix-halo/hsa_fix.sh new file mode 100755 index 0000000..a3ea9e0 --- /dev/null +++ b/docs/hardware/strix-halo/hsa_fix.sh @@ -0,0 +1,36 @@ +#!/usr/bin/env bash +# HSA runtime compatibility shim — Strix Halo (gfx1151) hardware profile — source this file to get hsa_fix_apply(). +# +# Official PyTorch +rocm7.0 wheels bundle a ROCm 7.0 HSA runtime that +# segfaults on gfx1151 (Strix Halo) when the kernel driver is from ROCm +# >= 7.1: crash in GpuAgent::QueueCreate while freezing the first code +# object. Preloading the SYSTEM runtime fixes it, because the system +# runtime always matches the installed kernel driver. +# +# Only applied when the system runtime is NEWER than the bundled one, so +# machines where the wheel's runtime works (or is newer) are untouched. +# +# Exports: LD_PRELOAD (extended when the fix applies) +# +# shellcheck disable=SC2034 # $2 (REPO_ROOT) kept for the sourced interface +hsa_fix_apply() { # $1 = venv python, $2 = repo root + local PY="$1" REPO_ROOT="$2" + local sys_rt="/opt/rocm/lib/libhsa-runtime64.so" + [ -f "$sys_rt" ] || return 0 + local sys_ver bundled_ver + sys_ver="$(cat /opt/rocm/.info/version 2>/dev/null || echo 0)" + bundled_ver="$("$PY" -c 'import torch; print(torch.version.hip or "0")' 2>/dev/null || echo 0)" + # numeric major.minor.patch comparison; preload only when system is newer + if awk -v a="$sys_ver" -v b="$bundled_ver" 'BEGIN { + split(a, A, "[^0-9]"); split(b, B, "[^0-9]"); + for (i = 1; i <= 3; i++) { if (A[i] + 0 != B[i] + 0) exit (A[i] + 0 > B[i] + 0) ? 0 : 1 } + exit 1 + }'; then + case ":${LD_PRELOAD:-}:" in + *":$sys_rt:"*) ;; # already included + *) export LD_PRELOAD="$sys_rt${LD_PRELOAD:+:${LD_PRELOAD}}" + echo "[sensenova-rocm] applied HSA runtime fix: LD_PRELOAD=$sys_rt (system ROCm ${sys_ver} > bundled ${bundled_ver})" + ;; + esac + fi +} diff --git a/docs/porting.md b/docs/porting.md new file mode 100644 index 0000000..0cf0cd1 --- /dev/null +++ b/docs/porting.md @@ -0,0 +1,140 @@ +# Porting Notes: SenseNova-U1.5 on AMD GPUs (ROCm) + +How the umbrella repo runs SenseNova-U1.5-8B-MoT on AMD GPUs, and why it +stays thin. Hardware-specific depth lives in +[docs/hardware/](hardware/) profiles; evidence-first claims live in +[docs/results/](results/). + +**TL;DR: the upstream code is already ROCm-ready** (one hardware-independent +interleave hotfix aside — §6). The NEO-unify reference implementation is +written against the `torch.cuda` namespace (which ROCm PyTorch provides) and +falls back to PyTorch SDPA when flash-attn is absent. The real work is +(1) getting PyTorch ROCm wheels that contain kernels for your gfx target, +and (2) fitting a ~50 GB bf16 checkpoint into your GPU's memory budget. +That's what `scripts/01-setup-venv.sh` and the `--vram_mode` flag take +care of. + +## 1. What "ROCm support" means in torch + +ROCm PyTorch wheels re-expose the GPU as the `cuda` device type: + +- `torch.cuda.is_available()` → `True` +- `device="cuda"` / `.to("cuda")` → runs on the AMD GPU +- `torch.version.hip` → the ROCm version the wheel was built against + +So any code that sticks to the `torch.cuda` API (like SenseNova-U1's +`src/sensenova_u1/utils/accel.py`, whose docstring explicitly says +"CUDA (incl. ROCm via the cuda namespace)") works unmodified. + +The one thing that does **not** carry over automatically is compiled +extension coverage: each wheel only contains GPU kernels for the gfx +architectures it was built for. `scripts/rocm_check.py` verifies your wheel +actually has kernels for your GPU before you waste an hour debugging a +`HIP error` deep inside model loading. + +## 2. Attention: SDPA instead of flash-attn + +Upstream prefers flash-attn (`--attn_backend auto` → flash if importable). +There is no flash-attn build for most ROCm gfx targets (including gfx1151). +Upstream's `auto` therefore selects **PyTorch SDPA** on ROCm, and this +repo's run scripts pin `--attn_backend sdpa` explicitly for determinism of +behavior. + +On ROCm, SDPA dispatches to an AOTriton flash kernel, an MIOpen fused +kernel, or the math fallback — all correct bf16 attention for the shapes +this model uses. + +Backend selection and determinism notes: see the per-hardware profile pages. + +## 3. Fitting a 50 GB checkpoint into smaller memories + +The bf16 checkpoint is ~50 GB. Two complementary mechanisms are used: + +### a) Unified-memory GTT spill (APUs) + +Unified-memory GTT spill on APUs: see +[docs/hardware/strix-halo/README.md](hardware/strix-halo/README.md). + +### b) Layer offload tiers (any GPU) + +Upstream ships a tiered layer-offload engine +(`sensenova_u1/utils/layer_offload.py`) exposed as `--vram_mode`: + +| vram_mode | strategy | suitable GPU memory | +|---|---|---| +| `full` | everything GPU-resident | ~52+ GiB discrete VRAM; on APUs the VRAM+GTT pool counts — see the [strix-halo profile](hardware/strix-halo/README.md) | +| `fast` | most layers resident; generation layers swapped per step | ~24 GB | +| `balanced` | heavier offload with async prefetch | ~16-24 GB | +| `low` | most aggressive offload | ~12-16 GB | + +In this repo `scripts/run-task.sh` injects `--vram_mode` for you and +defaults to `balanced`; select another tier with the `VRAM_MODE` +environment variable or by passing `--vram_mode` yourself. + +## 4. PyTorch wheel matrix (known ROCm wheel sources) + +| source | index | covers | +|---|---|---| +| **AMD official multi-arch** | `https://repo.amd.com/rocm/whl-multi-arch/` (`torch 2.12.0+rocm7.14.0`) | multi-gfx family wheels; gfx1100 and gfx1151 covered — gfx1151 runs without the HSA preload (verified on: gfx1151; see the [strix-halo profile](hardware/strix-halo/README.md)) | +| Official PyTorch ROCm 7 | `https://download.pytorch.org/whl/rocm7.0` (`torch 2.10.0+rocm7.0`) | gfx942, gfx950, gfx1100/1101, gfx1151, gfx1200 — gfx1151 needs the HSA preload (see the [strix-halo profile](hardware/strix-halo/README.md)) | +| AMD arch-specific nightlies | `https://rocm.nightlies.amd.com/v2//` | fallback for arches the official wheels miss | + +`scripts/01-setup-venv.sh` installs from a single index — `PY_INDEX_ROCM` +overrides; the default is the official PyTorch rocm6.3 index (`torch +2.8.0+rocm6.3`, the stack validated on gfx1100 in this repo) — and ends +with a GPU smoke test. `scripts/rocm_check.py` is the deeper check that +your wheel carries kernels for your gfx target. Validated stacks are +recorded with receipts under [docs/results/](results/). + +## 5. A real pitfall: the wheel's bundled HSA runtime (gfx1151) + +This pitfall is gfx1151-specific — see +[docs/hardware/strix-halo/README.md](hardware/strix-halo/README.md). + +## 6. The one upstream change: interleave hotfix + +Everything in the model code runs unmodified **except one upstream bugfix +touching 10 `_t2i_predict_v` call sites** we carry as a patch +([patches/0001-interleave-pass-image_size-to-_t2i_predict_v.patch](../patches/0001-interleave-pass-image_size-to-_t2i_predict_v.patch)): + +- `_t2i_predict_v()` dereferences `image_size` when the pixel head is + enabled (`use_pixel_head=true` in the released checkpoint), but the + interleave generation loops (`interleave_gen`, + `interleave_gen_image_only`) never pass it → + `TypeError: 'NoneType' object is not subscriptable` on the first + generated image. `t2i_generate` already passes it correctly. +- The bug is **hardware-independent** (the same code path exists on CUDA; + reproduced independently on gfx1100 and gfx1151) — see upstream + [PR #260](https://github.com/OpenSenseNova/SenseNova-U1/pull/260) + (open; local copy: + [docs/upstream/pr-interleave-image-size.md](upstream/pr-interleave-image-size.md)), + which also carries an independent cross-platform confirmation. This + repo's 0001 is regenerated to that PR's head — the reviewed version that + passes the per-image `cur_image_size` inside `interleave_gen_image_only`. +- `scripts/01-setup-venv.sh` applies the patch automatically after cloning + the pinned upstream checkout (`git apply`, idempotent — see + [patches/README.md](../patches/README.md)). + +If upstream merges the fix, delete the patch file and the setup step. + +## 7. What this repository changes vs upstream + +Almost nothing in the model code: three small patches, all documented in +[patches/README.md](../patches/README.md) — the interleave hotfix of §6 +(0001), an SDPA math-backend fallback for installs whose fused-SDPA wheels +are broken (0002, gfx1100 leaf wheels), and an opt-in +torch.compile/cudagraph acceleration (0003, off by default). +`scripts/01-setup-venv.sh`: + +1. installs ROCm PyTorch itself (default: the official PyTorch rocm6.3 + index; `PY_INDEX_ROCM` overrides — see §4) — upstream pins + `torch==2.8.0` on a cu128 index, an NVIDIA build, bypassed with + `pip install -e ... --no-deps`; +2. installs upstream from git at a **pinned commit** for reproducibility + and applies the patches; +3. wraps the four upstream example CLIs (`t2i`, `editing`, `vqa`, + `interleave`) with ROCm-verified defaults + ([scripts/run-task.sh](../scripts/run-task.sh)). + +If upstream ever ships ROCm wheels of its own, this repo shrinks to just +the verification harness — by design. diff --git a/docs/porting_CN.md b/docs/porting_CN.md new file mode 100644 index 0000000..40355fb --- /dev/null +++ b/docs/porting_CN.md @@ -0,0 +1,128 @@ +# 移植笔记：SenseNova-U1.5 在 AMD GPU 上（ROCm） + +umbrella 仓库如何在 AMD GPU 上运行 SenseNova-U1.5-8B-MoT，以及它为什么 +保持很薄。硬件相关的深入内容在 [docs/hardware/](hardware/) 的各硬件 +profile 页；证据优先的结论在 [docs/results/](results/)。 + +**一句话结论：上游代码本身就是 ROCm 就绪的**（仅一处硬件无关的 +interleave 热修复例外——见 §6）。NEO-unify 参考实现基于 `torch.cuda` +命名空间（ROCm PyTorch 原生提供），且在 flash-attn 缺失时自动回退 +PyTorch SDPA。真正的工作只有两件：(1) 拿到包含你的 gfx 目标内核的 +PyTorch ROCm wheel；(2) 把 ~50 GB 的 bf16 权重装进你的 GPU 内存预算。 +这正是 `scripts/01-setup-venv.sh` 和 `--vram_mode` 参数所处理的。 + +## 1. "ROCm 支持"在 torch 里意味着什么 + +ROCm 版 PyTorch 把 GPU 重新暴露为 `cuda` 设备类型： + +- `torch.cuda.is_available()` → `True` +- `device="cuda"` / `.to("cuda")` → 跑在 AMD GPU 上 +- `torch.version.hip` → wheel 构建所用的 ROCm 版本 + +所以只使用 `torch.cuda` API 的代码（如 SenseNova-U1 的 +`src/sensenova_u1/utils/accel.py`，其文档字符串明确写着 +"CUDA (incl. ROCm via the cuda namespace)"）**无需任何修改**即可运行。 + +唯一不能自动继承的是编译内核覆盖：每个 wheel 只包含它编译时指定的 +gfx 架构内核。`scripts/rocm_check.py` 会在你浪费数小时调试模型深处的 +`HIP error` 之前，先验证 wheel 是否真的带有你 GPU 的内核。 + +## 2. 注意力：用 SDPA 替代 flash-attn + +上游优先使用 flash-attn（`--attn_backend auto` → 可导入即用）。多数 +ROCm gfx 目标（包括 gfx1151）没有 flash-attn 构建。于是上游的 `auto` +在 ROCm 上会选择 **PyTorch SDPA**；本仓库的运行脚本显式固定 +`--attn_backend sdpa`，保证行为可复现。 + +ROCm 上 SDPA 分发到 AOTriton flash、MIOpen 融合内核或 math 回退—— +对本模型用到的形状都能给出正确的 bf16 注意力。 + +后端选择与确定性说明：见各硬件 profile 页。 + +## 3. 50 GB 权重如何装进更小的内存 + +bf16 权重约 50 GB。两套互补机制： + +### a) 统一内存 GTT 溢出（APU） + +统一内存 GTT 溢出（APU）：见 +[docs/hardware/strix-halo/README.md](hardware/strix-halo/README.md)。 + +### b) 分层 layer offload（任何 GPU） + +上游自带分层 offload 引擎（`sensenova_u1/utils/layer_offload.py`）， +通过 `--vram_mode` 暴露： + +| vram_mode | 策略 | 适用 GPU 内存 | +|---|---|---| +| `full` | 全部驻留 GPU | 需 ~52+ GiB 独立显存；APU 上可计入 VRAM+GTT 池——见 [strix-halo profile](hardware/strix-halo/README.md) | +| `fast` | 多数层驻留；生成层按步换入 | ~24 GB | +| `balanced` | 更重的 offload + 异步预取 | ~16-24 GB | +| `low` | 最激进 offload | ~12-16 GB | + +本仓库的 `scripts/run-task.sh` 会自动注入 `--vram_mode`，默认 +`balanced`；用 `VRAM_MODE` 环境变量（或自行传 `--vram_mode`）切换档位。 + +## 4. PyTorch wheel 矩阵（已知 ROCm wheel 来源） + +| 来源 | 索引 | 覆盖 | +|---|---|---| +| **AMD 官方多架构** | `https://repo.amd.com/rocm/whl-multi-arch/`（`torch 2.12.0+rocm7.14.0`） | 多 gfx 家族 wheel；gfx1100 与 gfx1151 均在覆盖内——gfx1151 无需 HSA 预加载（verified on: gfx1151；见 [strix-halo profile](hardware/strix-halo/README.md)） | +| 官方 PyTorch ROCm 7 | `https://download.pytorch.org/whl/rocm7.0`（`torch 2.10.0+rocm7.0`） | gfx942、gfx950、gfx1100/1101、gfx1151、gfx1200 —— gfx1151 需 HSA 预加载（见 [strix-halo profile](hardware/strix-halo/README.md)） | +| AMD 按架构 nightly | `https://rocm.nightlies.amd.com/v2//` | 官方 wheel 缺架构时的回退 | + +`scripts/01-setup-venv.sh` 从单一索引安装（`PY_INDEX_ROCM` 可覆盖； +默认为官方 PyTorch rocm6.3 索引，即 `torch 2.8.0+rocm6.3`——本仓库在 +gfx1100 上验证的栈），安装后跑 GPU 冒烟测试。`scripts/rocm_check.py` +是更深入的检查，验证 wheel 是否带有你 gfx 目标的内核。已验证的栈连同 +凭据记录在 [docs/results/](results/)。 + +## 5. 一个真实的坑：wheel 自带的 HSA 运行时（gfx1151） + +该坑仅影响 gfx1151——见 +[docs/hardware/strix-halo/README.md](hardware/strix-halo/README.md)。 + +## 6. 唯一的上游修改：interleave 热修复 + +模型代码**近乎零修改**，仅一处上游 bug 修复例外——改动 10 个 +`_t2i_predict_v` 调用点，以补丁形式携带 +（[patches/0001-interleave-pass-image_size-to-_t2i_predict_v.patch](../patches/0001-interleave-pass-image_size-to-_t2i_predict_v.patch)）： + +- 像素头启用时（发布权重 `use_pixel_head=true`）`_t2i_predict_v()` 会 + 解引用 `image_size`，但 interleave 生成循环（`interleave_gen`、 + `interleave_gen_image_only`）从不传它 → 首张生成图前即 + `TypeError: 'NoneType' object is not subscriptable`。`t2i_generate` + 一直传参正确。 +- 该 bug **与硬件无关**（CUDA 上同样存在这条代码路径；在 gfx1100 与 + gfx1151 上独立复现）——见上游 + [PR #260](https://github.com/OpenSenseNova/SenseNova-U1/pull/260) + （开放中；本地留存： + [docs/upstream/pr-interleave-image-size.md](upstream/pr-interleave-image-size.md)）， + 其中也带有一次独立跨平台确认。本仓库的 0001 按该 PR 的 head 重新 + 生成——即评审后的版本，在 `interleave_gen_image_only` 内改传每张图 + 的 `cur_image_size`。 +- `scripts/01-setup-venv.sh` 在克隆固定 commit 的上游后自动应用补丁 + （`git apply`，幂等——见 + [patches/README.md](../patches/README.md)）。 + +若上游合并该修复，删掉补丁文件与对应安装步骤即可。 + +## 7. 本仓库相对上游改了什么 + +模型代码：**近乎零修改**——共三个小补丁，全部记录在 +[patches/README.md](../patches/README.md)：§6 的 interleave 热修复 +（0001）、面向融合 SDPA wheel 损坏的安装的 math 回退（0002，gfx1100 +叶 wheel）、可选的 torch.compile/cudagraph 加速（0003，默认关闭）。 +`scripts/01-setup-venv.sh` 做的是： + +1. 自行安装 ROCm PyTorch（默认官方 PyTorch rocm6.3 索引， + `PY_INDEX_ROCM` 可覆盖——见 §4）——上游钉死 `torch==2.8.0` 于 + cu128 索引，是 NVIDIA 构建，用 `pip install -e ... --no-deps` + 绕开该钉死； +2. 从 git 以**固定 commit** 安装上游，保证可复现，并应用上述补丁； +3. 用 ROCm 验证过的默认值封装上游四个示例 CLI（`t2i`、`editing`、 + `vqa`、`interleave`，见 + [scripts/run-task.sh](../scripts/run-task.sh)）。 + +若上游某天原生提供 ROCm wheel，本仓库就收缩为纯验证套件——这是刻意 +设计。 diff --git a/docs/superpowers/plans/2026-08-26-umbrella-backport-8060s-wave1.md b/docs/superpowers/plans/2026-08-26-umbrella-backport-8060s-wave1.md new file mode 100644 index 0000000..5d20663 --- /dev/null +++ b/docs/superpowers/plans/2026-08-26-umbrella-backport-8060s-wave1.md @@ -0,0 +1,804 @@ +# Umbrella Backport 8060S — Wave 1 Implementation Plan + +> **For agentic workers:** REQUIRED SUB-SKILL: Use superpowers:subagent-driven-development (recommended) or superpowers:executing-plans to implement this plan task-by-task. Steps use checkbox (`- [ ]`) syntax for tracking. + +**Goal:** Execute wave-1 of the 8060S→umbrella backport: port the torch.compile patch + GPU tools + porting docs + Strix Halo hardware profile + upstream-tracking entry into this repo, close the governance gaps (CoC, SECURITY, templates, Discussions-ready CONTRIBUTING), and commit the alignment roadmap on the 8060S side. + +**Architecture:** All work happens in two local checkouts: the umbrella repo (this repo, at `/home/amd/Desktop/sensnova-merge/main-repo`, branch `backport/8060s-wave1`) and the 8060S repo (at `/home/amd/Desktop/SenseNova-U1.5-ROCm`, real git history + working venv, branch `docs/upstream-alignment`). Ports carry `Ported from` trailers; every task ends green on the repo's offline gates (`pytest tests/ -q`, `shellcheck`). No push happens in this plan — the final task prepares the PR and stops for user confirmation. + +**Tech Stack:** bash (strict mode, shellcheck-clean), stdlib-only Python 3, pytest offline (CPU-only, no GPU/network in CI), git plain-diff patches applied by the idempotent loop in `scripts/01-setup-venv.sh`. + +**Spec:** `docs/superpowers/specs/2026-08-26-umbrella-backport-8060s-design.md` — read it first; this plan argues from it. + +## Global Constraints + +- Repo layout gates: CI runs `python3 -m pytest tests/ -v` and `shellcheck scripts/*.sh scripts/lib/common.sh` — every task must keep both green offline (CI runners have no GPU, no torch, no network). +- Python ported into `scripts/` is stdlib-only and must not import torch at module level (CI has no torch; `rocm_check.py` imports torch only inside `main()` — preserve that). +- Patches in `patches/` are plain `git apply`-able diffs against upstream pin `76c32c2` (`scripts/lib/common.sh:14`); the setup loop auto-applies every `patches/*.patch`, so file 0003 needs no script change. +- Attribution: every ported commit message ends with a trailer line `Ported from AIwork4me/SenseNova-U1.5-ROCm-8060S@`. +- Bilingual discipline: every ported **doc** ships EN + CN in the same commit. +- Evidence labels: ported claims that are only gfx1151-verified carry `verified on: gfx1151 (Strix Halo)` plus a link into the 8060S repo until a gfx1100 receipt exists. +- Commit style: lowercase conventional prefixes used by this repo's history (`feat:`, `docs:`, `fix:`, `changelog:`). +- Git identity for commits: use `-c user.name=... -c user.email=...` only if repo config lacks identity; otherwise the configured identity. +- **Network constraint:** git transport to github.com is currently broken from this host (TLS resets). Work stays local; push/PR is a post-plan, user-confirmed step. + +## Environment Map (verify before starting) + +| what | path | state | +|---|---|---| +| umbrella repo (this repo) | `/home/amd/Desktop/sensnova-merge/main-repo` | tarball baseline `df59e51` + spec commit `b6c36a5` on `main`; **no real origin history** (graft before push, Task 10) | +| 8060S repo (port source, real history @ `28e4f3a`) | `/home/amd/Desktop/SenseNova-U1.5-ROCm` | clean tree expected; has `.venv` (torch 2.12.0+rocm7.14.0) | +| clean upstream checkout @ `76c32c2` | `/home/amd/SenseNova-U1.5-ROCm-work/upstream` | clean; use for `apply --check` only — never modify | +| gfx hardware | local iGPU gfx1151 (Strix Halo) | the 8060S reference host | + +Attribution SHAs (8060S history): tools + porting doc initial release `6046587f`; torch.compile patch `ceefff1a`; AOTriton#54 chain `d0af6efc`, `987015a1`, `28e4f3af`. + +--- + +### Task 1: Branch + patch 0003 (torch.compile/cudagraph safety) + +**Files:** +- Create: `patches/0003-optional-torch-compile-and-cudagraph-safety.patch` +- Modify: `patches/README.md` (append 0003 section) +- Modify: `CHANGELOG.md` (Unreleased → Added) + +**Interfaces:** +- Consumes: 8060S source `/home/amd/Desktop/SenseNova-U1.5-ROCm/patches/0002-optional-torch-compile-and-cudagraph-safety.patch` (60-line plain diff). +- Produces: `patches/0003-*.patch` auto-applied by `scripts/01-setup-venv.sh`'s glob loop; env knobs `SENSENOVA_COMPILE`, `SENSENOVA_COMPILE_SUPPRESS`, `SENSENOVA_COMPILE_FEATURES` documented in patches/README.md. + +- [ ] **Step 1: Create the wave-1 branch in the umbrella repo** + +```bash +cd /home/amd/Desktop/sensnova-merge/main-repo +git checkout -b backport/8060s-wave1 +``` + +Expected: `Switched to a new branch` (carries the spec commit `b6c36a5`). + +- [ ] **Step 2: Copy the patch with renumbered filename; content ports verbatim** + +```bash +cp /home/amd/Desktop/SenseNova-U1.5-ROCm/patches/0002-optional-torch-compile-and-cudagraph-safety.patch \ + patches/0003-optional-torch-compile-and-cudagraph-safety.patch +``` + +- [ ] **Step 3: Verify it applies cleanly to the pinned upstream (no tree modification)** + +```bash +git -C /home/amd/SenseNova-U1.5-ROCm-work/upstream apply --check \ + patches/0003-optional-torch-compile-and-cudagraph-safety.patch && echo APPLY-OK +``` + +Expected: `APPLY-OK`. If it fails: copy the upstream checkout to `/tmp/upstream-scratch`, `git -C /tmp/upstream-scratch apply --3way` the patch, fix hunks, regenerate with `git -C /tmp/upstream-scratch diff > patches/0003-…patch`, and record the regeneration in the commit message. + +- [ ] **Step 4: Runtime smoke on gfx1151 (the patch's original verification host)** + +```bash +cd /home/amd/Desktop/SenseNova-U1.5-ROCm +SENSENOVA_COMPILE=1 SENSENOVA_SIZE=1024x1024 SENSENOVA_STEPS=4 SENSENOVA_SEED=42 \ +bash scripts/run-t2i.sh "a red lighthouse on a rocky cliff, smoke test" /tmp/wave1-compile-smoke.png +``` + +Expected: log line `[compile] enabled {'dynamic': False}` and a saved PNG. (This runs the 8060S venv, whose installed tree carries the identical change — Step 3 proved our file is the same content against the same pin.) If the model cache/venv is unavailable, skip with a note in the PR description and cite the existing V1–V8 evidence in the 8060S repo. + +- [ ] **Step 5: Append the 0003 section to `patches/README.md`** + +Add after the 0002 section (match its heading style): + +```markdown +## 0003-optional-torch-compile-and-cudagraph-safety.patch + +**Opt-in acceleration (all platforms).** Adds env-gated `torch.compile` +around `_t2i_predict_v` (and optionally `extract_feature`/`patchify`): +`SENSENOVA_COMPILE=1|default|reduce-overhead|…` enables, +`SENSENOVA_COMPILE_SUPPRESS=1` sets dynamo suppress-errors, +`SENSENOVA_COMPILE_FEATURES=1` also compiles the feature/patchify path. +The model-side hunk marks cudagraph step boundaries and materializes the +conditioned prediction before the unconditioned re-run of the same +compiled callable — without it, cudagraph output pooling overwrites the +CFG condition with the uncondition result. + +verified on: gfx1151 (Strix Halo) — V1–V8 speedup + Qwen-Image-Bench +paired quality validation live in the +[8060S repo](https://github.com/AIwork4me/SenseNova-U1.5-ROCm-8060S) +(`evidence/speedup2/`, PASS-with-caveat: compile stays opt-in until a +gfx1100 receipt exists here). +``` + +- [ ] **Step 6: CHANGELOG Unreleased entry** + +Add under `## [Unreleased]` → `### Added`: + +```markdown +- Patch 0003: optional torch.compile/cudagraph safety (env-gated + `SENSENOVA_COMPILE*`); ported from the Strix Halo sibling repo — + verified on gfx1151 (V1–V8 + quality bench there); gfx1100 receipt + pending, compile stays opt-in. +``` + +- [ ] **Step 7: Commit** + +```bash +cd /home/amd/Desktop/sensnova-merge/main-repo +git add patches/0003-optional-torch-compile-and-cudagraph-safety.patch patches/README.md CHANGELOG.md +git commit -m "feat: patch 0003 — optional torch.compile/cudagraph safety (env-gated) + +Ported from AIwork4me/SenseNova-U1.5-ROCm-8060S@ceefff1a" +``` + +--- + +### Task 2: Port `gpu_monitor.py` (VRAM+GTT sampling) + +**Files:** +- Create: `scripts/gpu_monitor.py` +- Test: `tests/test_scripts.py` (append `test_gpu_monitor_*`) + +**Interfaces:** +- Consumes: 8060S source `/home/amd/Desktop/SenseNova-U1.5-ROCm/python/gpu_monitor.py` (100 lines, stdlib-only). +- Produces: CLI `scripts/gpu_monitor.py --interval --out `; JSONL lines `{"ts":…, "vram_used_gib":…, "gtt_used_gib":…}` + final `{"summary":true, "samples":n, "peak_vram_gib":…, "peak_gtt_gib":…, "peak_total_gib":…}`; importable functions `_used_bytes(which:str)->int|None`, `sample_rocm_smi()->tuple[float,float]|None`. + +- [ ] **Step 1: Write the failing tests (append to `tests/test_scripts.py`)** + +```python +def test_gpu_monitor_help(): + r = subprocess.run( + [sys.executable, f"{ROOT}/scripts/gpu_monitor.py", "--help"], + capture_output=True, text=True) + assert r.returncode == 0, r.stderr + assert "VRAM" in r.stdout and "GTT" in r.stdout + + +def test_gpu_monitor_parses_rocm_smi_csv(): + import importlib.util + spec = importlib.util.spec_from_file_location( + "gpu_monitor", f"{ROOT}/scripts/gpu_monitor.py") + gm = importlib.util.module_from_spec(spec) + spec.loader.exec_module(gm) + + fake_csv = ("device,vram total (B),vram used (B)\n" + "card0,34359738368,10737418240\n" + "card1,34359738368,2147483648\n") + class FakeProc: + returncode = 0 + stdout = fake_csv + real_run = gm.subprocess.run + gm.subprocess.run = lambda *a, **k: FakeProc() + try: + assert gm._used_bytes("vram") == 10737418240 # max across cards + vram, gtt = gm.sample_rocm_smi() + assert abs(vram - 10.0) < 0.01 and abs(gtt - 10.0) < 0.01 # /2**30 + finally: + gm.subprocess.run = real_run +``` + +(No pytest fixtures: this file has a `__main__` runner that calls every `test_*` with no arguments — new tests must stay fixture-free.) + +- [ ] **Step 2: Run to verify failure** + +Run: `cd /home/amd/Desktop/sensnova-merge/main-repo && python3 -m pytest tests/ -q -k gpu_monitor` +Expected: FAIL (file not found / import error). + +- [ ] **Step 3: Port the file with two doc adaptations** + +```bash +cp /home/amd/Desktop/SenseNova-U1.5-ROCm/python/gpu_monitor.py scripts/gpu_monitor.py +``` + +Then edit `scripts/gpu_monitor.py` docstring (code body unchanged): +- Usage line: `python gpu_monitor.py --interval 0.5 --out evidence/mem_usage.jsonl` → `python scripts/gpu_monitor.py --interval 0.5 --out docs/results/logs/mem_usage.jsonl` +- Last docstring line: `…aggregation by scripts/verify-all.sh.` → `…aggregation by scripts/validate.sh receipts.` + +- [ ] **Step 4: Run tests to verify pass + full suite green** + +Run: `python3 -m pytest tests/ -q` +Expected: PASS (all, including pre-existing). + +- [ ] **Step 5: Real-GPU sanity on gfx1151 (not a CI gate)** + +```bash +timeout 5 python3 scripts/gpu_monitor.py --interval 0.5 --out /tmp/mem.jsonl; tail -1 /tmp/mem.jsonl +``` + +Expected: a summary JSON with nonzero `peak_vram_gib` (timeout exit 124 is fine — summary prints only on SIGINT/SIGTERM, so check sampled lines: `head -2 /tmp/mem.jsonl`). + +- [ ] **Step 6: CHANGELOG + commit** + +Add to Unreleased → Added: + +```markdown +- `scripts/gpu_monitor.py`: sample VRAM + GTT while a workload runs + (JSONL + peak summary) — GTT visibility is what evidences APU + GTT-spill runs. +``` + +```bash +git add scripts/gpu_monitor.py tests/test_scripts.py CHANGELOG.md +git commit -m "feat: port gpu_monitor.py (VRAM+GTT sampling) from 8060S + +Ported from AIwork4me/SenseNova-U1.5-ROCm-8060S@6046587f" +``` + +--- + +### Task 3: Port `rocm_check.py` (env sanity, GPU-dependent parts stay runtime-only) + +**Files:** +- Create: `scripts/rocm_check.py` +- Test: `tests/test_scripts.py` (append `test_rocm_check_*`) + +**Interfaces:** +- Consumes: 8060S source `/home/amd/Desktop/SenseNova-U1.5-ROCm/python/rocm_check.py` (171 lines; imports torch ONLY inside `main()` — keep it that way). +- Produces: CLI `scripts/rocm_check.py [--alloc-test GIB] [--json]`; exit 0 = env OK; `finish(report:dict, as_json:bool, ok:bool)->int` importable offline. + +- [ ] **Step 1: Write the failing tests** + +```python +def test_rocm_check_help(): + r = subprocess.run( + [sys.executable, f"{ROOT}/scripts/rocm_check.py", "--help"], + capture_output=True, text=True) + assert r.returncode == 0, r.stderr + assert "alloc" in r.stdout.lower() + + +def test_rocm_check_finish_formats_report(): + import importlib.util, io, contextlib + spec = importlib.util.spec_from_file_location( + "rocm_check", f"{ROOT}/scripts/rocm_check.py") + rc = importlib.util.module_from_spec(spec) + spec.loader.exec_module(rc) + buf = io.StringIO() + with contextlib.redirect_stdout(buf): + code = rc.finish({"torch": "x", "ok": False}, as_json=False, ok=False) + assert code == 1 and "[FAIL]" in buf.getvalue() +``` + +- [ ] **Step 2: Run to verify failure** + +Run: `python3 -m pytest tests/ -q -k rocm_check` +Expected: FAIL (file not found). + +- [ ] **Step 3: Port with doc adaptations** + +```bash +cp /home/amd/Desktop/SenseNova-U1.5-ROCm/python/rocm_check.py scripts/rocm_check.py +``` + +Docstring edits only (body unchanged): +- `(see scripts/setup-rocm.sh)` → `(see scripts/01-setup-venv.sh)` +- Line 5-8 numbered list stays; line about `--alloc-test` keeps the APU/GTT-spill explanation (that is the strix-halo profile hook). + +- [ ] **Step 4: Run full suite** + +Run: `python3 -m pytest tests/ -q` +Expected: PASS. (If a test imports torch accidentally it fails on CI-like hosts — the module import in Step 1 must stay torch-free; it is, torch is inside `main()`.) + +- [ ] **Step 5: Real run on gfx1151 (not a CI gate)** + +```bash +python3 scripts/rocm_check.py --alloc-test 8 --json | head -20 +``` + +Expected: `"ok": true`, `alloc_test: "ok"` (8 GiB fits VRAM carve-out; a `--alloc-test 48` run belongs to the strix-halo profile docs, not here). + +- [ ] **Step 6: CHANGELOG + commit** + +Unreleased → Added: + +```markdown +- `scripts/rocm_check.py`: env sanity check (HIP build, wheel gfx + coverage, bf16 matmul + SDPA smoke, optional GTT-spill alloc test). +``` + +```bash +git add scripts/rocm_check.py tests/test_scripts.py CHANGELOG.md +git commit -m "feat: port rocm_check.py (environment sanity) from 8060S + +Ported from AIwork4me/SenseNova-U1.5-ROCm-8060S@6046587f" +``` + +--- + +### Task 4: `receipt.py` hardware field (test + doc only — code is already generic) + +**Files:** +- Modify: `tests/test_scripts.py` (`test_receipt_writer` extension) +- Modify: `scripts/receipt.py` (docstring line only) +- Modify: `CONTRIBUTING.md` (evidence-policy note — full section lands in Task 8) + +**Interfaces:** +- Consumes: `receipt.py` already stores arbitrary `key=value` pairs (`parse_value` tries JSON, else string) — no code change. +- Produces: documented convention `hardware=gfx1100|gfx1151|…` in receipts; asserted round-trip in tests. + +- [ ] **Step 1: Extend the failing test — in `test_receipt_writer`, change the receipt invocation and assertions** + +Replace the `subprocess.run` arg list entry `"note=plain text"` args with: + +```python + [sys.executable, f"{ROOT}/scripts/receipt.py", out, + "block=t2i", "wall_seconds=12.5", "hardware=gfx1151", + f"sha256:{art}", "note=plain text"], +``` + +and add after `assert receipt["note"] == "plain text"`: + +```python + assert receipt["hardware"] == "gfx1151" +``` + +- [ ] **Step 2: Run — this PASSES immediately (proves genericity); that is the point** + +Run: `python3 -m pytest tests/ -q -k receipt` +Expected: PASS. If it fails, `receipt.py` regressed — fix before proceeding. + +- [ ] **Step 3: Document the field in `scripts/receipt.py` docstring** + +Add one line to the module docstring after the "Special keys" paragraph: + +```markdown +Conventional keys recorded by validate.sh runs: `block`, `hardware` +(gfx arch the run executed on, e.g. gfx1100 / gfx1151 — lets +cross-hardware evidence be queried), `rocm_stack`. +``` + +- [ ] **Step 4: Commit** + +```bash +git add tests/test_scripts.py scripts/receipt.py +git commit -m "test+docs: receipt hardware field convention (code already generic)" +``` + +--- + +### Task 5: `docs/porting.md` + `docs/porting_CN.md` (generic split of ROCm_PORTING) + +**Files:** +- Create: `docs/porting.md` (EN), `docs/porting_CN.md` (CN) +- Modify: `docs/getting-started.md` (one cross-link line) + +**Interfaces:** +- Consumes: `/home/amd/Desktop/SenseNova-U1.5-ROCm/docs/ROCm_PORTING.md` (146 lines, EN) and `…/ROCm_PORTING_CN.md` (107 lines, CN). +- Produces: umbrella-generic porting notes; APU-specific sections are NOT here (they move in Task 6). + +- [ ] **Step 1: Write `docs/porting.md` from the EN source with this section mapping** + +| 8060S source section | destination | +|---|---| +| header + TL;DR | keep, reword "this repository" → umbrella repo; drop `setup-rocm.sh`/`vram_mode tiering automate` phrasing → `scripts/01-setup-venv.sh` + `--vram_mode` | +| §1 What "ROCm support" means | keep; `python/rocm_check.py` → `scripts/rocm_check.py` | +| §2 Attention: SDPA | keep through "math fallback"; DELETE the AOTriton-enablement + hd72 caveat paragraph (8060S-specific stack guidance) and the determinism notes paragraph (that is their PERFORMANCE.md content) — replace with one line: `Backend selection and determinism notes: see the per-hardware profile pages.` | +| §3 intro + §3b layer offload tiers | keep (generic, any GPU) | +| §3a GTT spill (APUs) | NOT here — one pointer line: `Unified-memory GTT spill on APUs: see docs/hardware/strix-halo/README.md.` | +| §4 wheel matrix | keep table shell, but generic: drop the "default since 2026-08-23 / 8060S A/B" framing; cover the official pytorch rocm7.0 index + AMD multi-arch + nightlies rows; gfx1151-specific HSA note → pointer to strix-halo profile | +| §5 HSA runtime pitfall | NOT here — pointer to strix-halo profile | +| §6 interleave hotfix | keep, but reference THIS repo's patch 0001 (same fix, our regenerated version) and upstream PR#260 | +| §7 what this repo changes | keep, adapted: `setup-rocm.sh` → `01-setup-venv.sh`; mention patches 0001–0003 (0003 = opt-in compile) | + +Header of the new file: + +```markdown +# Porting Notes: SenseNova-U1.5 on AMD GPUs (ROCm) + +How the umbrella repo runs SenseNova-U1.5-8B-MoT on AMD GPUs, and why it +stays thin. Hardware-specific depth lives in +[docs/hardware/](hardware/) profiles; evidence-first claims live in +[docs/results/](results/). +``` + +- [ ] **Step 2: Write `docs/porting_CN.md` by applying the same mapping to the CN source** (`ROCm_PORTING_CN.md`, on disk). Same deletions/pointers, translated consistently with the EN file's terminology. + +- [ ] **Step 3: Cross-link from `docs/getting-started.md`** + +Add near the top of getting-started.md (after its intro line): + +```markdown +For how this port works (wheels, attention backend, memory tiers): +[Porting Notes](porting.md) · [移植笔记](porting_CN.md) +``` + +- [ ] **Step 4: Commit** + +```bash +git add docs/porting.md docs/porting_CN.md docs/getting-started.md +git commit -m "docs: porting notes (EN+CN) — generic split from 8060S ROCm_PORTING + +Ported from AIwork4me/SenseNova-U1.5-ROCm-8060S@6046587f" +``` + +--- + +### Task 6: Strix Halo hardware profile (`docs/hardware/strix-halo/`) + +**Files:** +- Create: `docs/hardware/strix-halo/README.md` (EN), `docs/hardware/strix-halo/README_CN.md` (CN) +- Create: `docs/hardware/strix-halo/hsa_fix.sh` (ported) +- Modify: `.github/workflows/ci.yml` (shellcheck the new script) +- Modify: `CHANGELOG.md` + +**Interfaces:** +- Consumes: `/home/amd/Desktop/SenseNova-U1.5-ROCm/scripts/hsa_fix.sh` (34 lines, shell function `hsa_fix_apply`); 8060S porting doc §3a (GTT spill) + §5 (HSA pitfall) content. +- Produces: profile dir referenced by `docs/porting.md` (Task 5 pointer); `hsa_fix_apply "$PY" "$REPO_ROOT"` shell function, sourced not executed. + +- [ ] **Step 1: Port `hsa_fix.sh` with header rewording** + +```bash +mkdir -p docs/hardware/strix-halo +cp /home/amd/Desktop/SenseNova-U1.5-ROCm/scripts/hsa_fix.sh docs/hardware/strix-halo/hsa_fix.sh +``` + +Edit the header comment block (function body unchanged): replace `sourced by setup-rocm.sh and common.sh` with `Strix Halo (gfx1151) hardware profile — source this file to get hsa_fix_apply()`; keep the mechanism comment intact. + +- [ ] **Step 2: Extend CI shellcheck to cover it** + +In `.github/workflows/ci.yml`, change: + +```yaml + shellcheck scripts/*.sh scripts/lib/common.sh +``` + +to: + +```yaml + shellcheck scripts/*.sh scripts/lib/common.sh docs/hardware/strix-halo/hsa_fix.sh +``` + +Local gate: `shellcheck docs/hardware/strix-halo/hsa_fix.sh` → no output (clean). + +- [ ] **Step 3: Write `docs/hardware/strix-halo/README.md`** + +Structure (content adapted from 8060S porting §3a + §5 + their PERFORMANCE/INSTALL facts on disk at `/home/amd/Desktop/SenseNova-U1.5-ROCm/docs/`): + +```markdown +# Hardware Profile: Strix Halo (Radeon 8060S iGPU, gfx1151) + +verified on: gfx1151 — Ryzen AI MAX+ PRO 395, ROCm 7.2.1 system + +torch 2.12.0+rocm7.14.0 wheels. Full evidence corpus (V1–V8 speedup, +quality bench, determinism, tier tests): +[SenseNova-U1.5-ROCm-8060S](https://github.com/AIwork4me/SenseNova-U1.5-ROCm-8060S). +gfx1100 receipts live in [docs/results/](../../results/) here. + +## Why full bf16 fits a "32 GB" iGPU — GTT spill +<§3a content: VRAM carve-out vs GTT pools, same-package LPDDR5X, allocator +spill, `scripts/rocm_check.py --alloc-test 48` as the verifier> + +## HSA runtime pitfall (torch +rocm7.0 wheels, driver ≥ 7.1) +<§5 content: GpuAgent::QueueCreate segfault, system-runtime preload, +hsa_fix.sh applies it only when the system runtime is newer> + +## Wheel guidance for gfx1151 + + +## AOTriton caveat (hd=72 ViT heads) + + +## Companion repo +daily-driver scripts, one-command quickstart, and the full evidence +corpus live in the 8060S repo (the Strix Halo lab); experiments there +graduate into this umbrella via PR when they meet the graduation +criteria (CONTRIBUTING.md). +``` + +- [ ] **Step 4: Write `README_CN.md` (same structure, CN terminology consistent with 8060S `ROCm_PORTING_CN.md`)** + +- [ ] **Step 5: CHANGELOG + commit** + +Unreleased → Added: + +```markdown +- Hardware profile `docs/hardware/strix-halo/` (EN+CN): GTT-spill + memory model, HSA runtime shim (`hsa_fix.sh`, now CI-shellchecked), + gfx1151 wheel guidance, AOTriton hd72 caveat — the umbrella's first + APU profile, ported from the 8060S repo. +``` + +```bash +git add docs/hardware .github/workflows/ci.yml CHANGELOG.md +git commit -m "feat: strix-halo hardware profile (EN+CN docs + hsa_fix shim, CI-covered) + +Ported from AIwork4me/SenseNova-U1.5-ROCm-8060S@6046587f" +``` + +--- + +### Task 7: `docs/upstream/aotriton-54.md` (upstream tracking entry) + +**Files:** +- Create: `docs/upstream/aotriton-54.md` + +**Interfaces:** +- Consumes: 8060S evidence chain at `/home/amd/Desktop/SenseNova-U1.5-ROCm/` — `evidence/` V2 notes, `docs/superpowers/specs/2026-08-23-speedup-and-qwen-image-bench-design.md` (its VERIFICATION notes), commits `d0af6efc`/`987015a1`/`28e4f3af`. +- Produces: standalone tracking doc parallel to the existing `docs/upstream/` filings (read one for format before writing). + +- [ ] **Step 1: Read an existing `docs/upstream/` file for format, then write the entry** + +Content requirements (facts to verify against the 8060S sources above — do not invent): affected component (AOTriton flash kernel, hd=72 head-dim ViT configs), symptom (launch failure / NaN drift observed in speedup2 re-verification), root-cause status (aotriton issue #54 filed 2026-08-26, fix branch unmerged), cross-verified on the 7.14 stack, link to the 8060S evidence directory, and the caveat chain for foreign models (hd≠64 configs until fixed upstream). + +- [ ] **Step 2: Commit** + +```bash +git add docs/upstream/aotriton-54.md +git commit -m "docs: upstream tracking — AOTriton hd72 kernel bug (issue #54) + +Ported from AIwork4me/SenseNova-U1.5-ROCm-8060S@d0af6efc 987015a1 28e4f3af" +``` + +--- + +### Task 8: Governance additions (CoC, SECURITY, templates, CONTRIBUTING, README) + +**Files:** +- Create: `CODE_OF_CONDUCT.md`, `SECURITY.md` +- Create: `.github/ISSUE_TEMPLATE/bug_report.md`, `.github/ISSUE_TEMPLATE/validation_report.md`, `.github/PULL_REQUEST_TEMPLATE.md` +- Modify: `CONTRIBUTING.md` (cross-repo section), `README.md` + `README_CN.md` (hardware coverage + sibling link) + +**Interfaces:** +- Consumes: CONTRIBUTING's existing "validation: " issue convention. +- Produces: governance surface a community expects; CONTRIBUTING §"Contributing between the two repos" referenced by the 8060S-side roadmap (Task 9). + +- [ ] **Step 1: `CODE_OF_CONDUCT.md`** — concise, bilingual note at end: + +```markdown +# Code of Conduct + +## Our pledge +Be evidence-first about claims and kind to people. We welcome +contributors of every background and skill level. + +## Standards +**Encouraged:** citing receipts for claims, asking questions, reporting +both successes and failures, reviewing others' work respectfully. +**Unacceptable:** harassment or discriminatory language, hostile +responses to newcomers, presenting unverified numbers as measured, +doxxing, spam. + +## Enforcement +Report violations by opening a private GitHub security advisory +("Report a vulnerability" → select "not applicable") or contacting a +maintainer via GitHub. Maintainers will respond within 7 days and may +remove comments/commits/ban repeat offenders. + +## Attribution +Adapted in spirit from the Contributor Covenant 2.1. +本行为准则同样适用于中文社区讨论；违反行为可同样通过上述渠道举报。 +``` + +- [ ] **Step 2: `SECURITY.md`** + +```markdown +# Security Policy + +## Supported versions +The tip of `main` only. + +## Reporting a vulnerability +Please do NOT open a public issue. Use GitHub's private vulnerability +reporting ("Security" tab → "Report a vulnerability"). Include +reproduction steps and affected components (scripts/, patches/, docs/ +pipeline). Expect a response within 7 days. + +## Scope +This repo ships shell/python scripts and upstream patches; it does not +serve anything to the network. Model weights come from upstream +ModelScope — weight integrity is verified by +configs/artifact-manifest.json SHA256s (any tampering report is +welcome but belongs upstream). +``` + +- [ ] **Step 3: Issue templates** — `.github/ISSUE_TEMPLATE/bug_report.md`: + +```markdown +--- +name: Bug report +about: Something broke on your AMD GPU +title: "bug: " +labels: bug +--- +**Hardware / stack** (paste `scripts/00-check-env.sh` output or +docs/results/environment.json): +**Exact command:** +**Full log** (pastebin/gist if long): +**What you expected vs what happened:** +``` + +`.github/ISSUE_TEMPLATE/validation_report.md`: + +```markdown +--- +name: Validation report +about: Share results from YOUR hardware (successes AND failures) +title: "validation: / ROCm " +labels: validation +--- +**GPU / ROCm / torch:** +**`00-check-env.sh` output:** +**Summary table** (`python3 scripts/summarize_results.py`): +**Receipts** (docs/results/… if you can share): +**Anything that failed, with the log:** +``` + +`.github/PULL_REQUEST_TEMPLATE.md`: + +```markdown +**What & why** (one paragraph; link the spec/design doc if any): + +**Evidence** (receipts under docs/results/, or `verified on: ` + +link for ports from sibling repos — see CONTRIBUTING): + +**Gates** (must both be green): +- [ ] `python3 -m pytest tests/ -q` +- [ ] `shellcheck scripts/*.sh scripts/lib/common.sh docs/hardware/strix-halo/hsa_fix.sh` + +**If ported from SenseNova-U1.5-ROCm-8060S:** source commit SHA: +``` + +- [ ] **Step 4: CONTRIBUTING.md cross-repo section** (append at end): + +```markdown +## Contributing between the two repos + +This repo is the **umbrella** (hardware-generic assets + hardware +profiles). [SenseNova-U1.5-ROCm-8060S](…) is the **Strix Halo lab**: +APU-first experiments live there; verified work graduates here. + +**8060S → umbrella (graduation).** A change graduates when it is +① hardware-independent or absorbed as a hardware profile, ② verified on +gfx1151 with evidence in the 8060S repo, ③ passes this repo's offline +CI, ④ documented EN+CN, ⑤ not a duplicate of an existing tool here. +Open a PR labelled `backport`; the commit carries +`Ported from …@` and the claim is labelled +`verified on: gfx1151` until a gfx1100 receipt lands here. + +**Umbrella → 8060S (flow-down).** Changes to shared surfaces +(`patches/`, `scripts/`) are cherry-picked by the 8060S repo as +needed; that repo records `Merged-upstream: `. + +**Evidence across repos.** `receipt.py` records `hardware=` on +every run; cross-hardware claims must link the producing repo's +evidence. +``` + +Fill the `…` link with `https://github.com/AIwork4me/SenseNova-U1.5-ROCm-8060S`. + +- [ ] **Step 5: README hardware coverage** — in `README.md`, after the intro paragraph (before the showcase section), insert: + +```markdown +## Hardware coverage + +| hardware | status | where | +|---|---|---| +| gfx1100 (RDNA3 dGPU) | fully validated — receipts in [docs/results/](docs/results/) | this repo | +| gfx1151 (Strix Halo APU) | fully validated — evidence in the [8060S repo](https://github.com/AIwork4me/SenseNova-U1.5-ROCm-8060S); profile: [docs/hardware/strix-halo/](docs/hardware/strix-halo/) | this repo + Strix Halo lab | + +Other AMD GPUs: the pipeline is hardware-generic (see +[Porting Notes](docs/porting.md)) — validation reports welcome +(issue template `validation_report`). +``` + +Mirror the same table in `README_CN.md` (translated, same links). Add the 8060S link line to README_CN's intro too. + +- [ ] **Step 6: Full gates + commit** + +Run: `python3 -m pytest tests/ -q && shellcheck scripts/*.sh scripts/lib/common.sh docs/hardware/strix-halo/hsa_fix.sh` +Expected: all green. + +```bash +git add CODE_OF_CONDUCT.md SECURITY.md .github CONTRIBUTING.md README.md README_CN.md +git commit -m "docs: governance — CoC, SECURITY, issue/PR templates, cross-repo contributing, hardware coverage" +``` + +--- + +### Task 9: 8060S-side alignment roadmap (docs-only, separate repo) + +**Files (in `/home/amd/Desktop/SenseNova-U1.5-ROCm`, branch `docs/upstream-alignment`):** +- Create: `docs/UPSTREAM.md` +- Modify: `README.md` (upstream note after the badges) + +**Interfaces:** +- Consumes: the spec §4 (three alignment steps) + §5 (graduation criteria — final wording lives in the umbrella CONTRIBUTING from Task 8; reference, don't fork). +- Produces: the 8060S-side commitment doc; no behavior change in that repo. + +- [ ] **Step 1: Branch + write `docs/UPSTREAM.md`** + +```bash +cd /home/amd/Desktop/SenseNova-U1.5-ROCm && git checkout -b docs/upstream-alignment +``` + +```markdown +# Upstream relationship & alignment roadmap + +Upstream umbrella: [SenseNova-U1.5-ROCm](https://github.com/AIwork4me/SenseNova-U1.5-ROCm). +This repo stays **independent** — its own quickstart, evidence corpus, +announcements, releases — while aligning structure so work flows both +ways cheaply. + +## Graduation (this repo → umbrella) +Experiments here graduate into the umbrella when they meet the +criteria in the umbrella's CONTRIBUTING ("Contributing between the two +repos"): hardware-independent or profile-absorbed, gfx1151 evidence, +CI-green, EN+CN docs, no duplicate tool. Graduating PRs carry +`Ported from …@`; this repo records `Merged-upstream: `. + +## Alignment roadmap (this repo, own pace) +1. **Governance**: CONTRIBUTING, CHANGELOG (backfilled), Code of + Conduct, SECURITY, issue/PR templates — adapted from the umbrella. +2. **Structure**: `python/*` tools adopt the umbrella's `scripts/` + names once wave-1 lands there (gpu_monitor, rocm_check); root + `quickstart.sh` stays as a thin wrapper. +3. **Quality gates**: the umbrella's CI (shellcheck strict + offline + pytest) and a `tests/` skeleton. + +## Flow-down (umbrella → this repo) +Umbrella changes to `patches/`/`scripts/` are cherry-picked here as +needed. +``` + +- [ ] **Step 2: README upstream note** — after the two badge lines add: + +```markdown +> Upstream umbrella (multi-hardware): [SenseNova-U1.5-ROCm](https://github.com/AIwork4me/SenseNova-U1.5-ROCm) — this repo is the independent Strix Halo lab; see [docs/UPSTREAM.md](docs/UPSTREAM.md). +``` + +Mirror in `README_CN.md` (same position, CN wording). + +- [ ] **Step 3: Commit (8060S repo)** + +```bash +git add docs/UPSTREAM.md README.md README_CN.md +git commit -m "docs: upstream relationship & alignment roadmap (umbrella graduation model)" +``` + +--- + +### Task 10: Final verification, PR preparation, STOP + +**Files:** +- Create: `/home/amd/Desktop/sensnova-merge/WAVE1-PR.md` (PR description, outside both repos) + +- [ ] **Step 1: Full offline gates in the umbrella repo** + +```bash +cd /home/amd/Desktop/sensnova-merge/main-repo +python3 -m pytest tests/ -v +shellcheck scripts/*.sh scripts/lib/common.sh docs/hardware/strix-halo/hsa_fix.sh +bash -n scripts/*.sh +``` + +Expected: pytest all green; shellcheck silent; bash -n silent. + +- [ ] **Step 2: Verify every commit carries its trailer** + +```bash +git log --format='%h %s%n%b' main..backport/8060s-wave1 | grep -B2 'Ported from' +``` + +Expected: 4 ported commits (Tasks 1, 2, 3, 5/6/7 share trailers) each showing the trailer. + +- [ ] **Step 3: Write `WAVE1-PR.md`** + +Must contain: summary paragraph (umbrella wave-1), per-commit list with SHA ranges, the `verified on: gfx1151` claims inventory (patch 0003, strix-halo profile) with 8060S evidence links, explicit "gfx1100 re-validation pending" note, and the graduation-criteria mapping (CONTRIBUTING §cross-repo). + +- [ ] **Step 4: STOP — user gate** + +Do NOT push. Report to the user: branch names, commit list, gate results, smoke-test outcomes, and the push procedure for when they confirm: + +```bash +# when git transport to github.com works again (TLS currently broken on this host): +cd /home/amd/Desktop/sensnova-merge/main-repo +git fetch origin && git rebase --onto origin/main backport/8060s-wave1 +git push -u origin backport/8060s-wave1 +# then open the PR from WAVE1-PR.md; likewise push 8060S docs/upstream-alignment +``` + +(`` = the commit before the first wave-1 commit; recorded at execution time.) + +--- + +## Self-Review (done at plan-writing time) + +- **Spec coverage:** inventory items 1–6 → Tasks 1–7; repo-side adaptations (receipt hardware field, governance, README) → Tasks 4+8; 8060S roadmap → Task 9; evidence/attribution rules → trailers + labels throughout; execution boundary/STOP → Task 10. Deferred items (bench_report, evidence bulk, gfx1100 re-validation, 8060S restructuring execution) match the spec's deferrals. +- **Type consistency:** `gpu_monitor._used_bytes/sample_rocm_smi`, `rocm_check.finish(report, as_json, ok)`, `hsa_fix_apply "$PY" "$REPO_ROOT"` — names match source files verbatim. +- **Placeholders:** none — every new file's content is embedded or mapped section-by-section from an on-disk source; no "TBD"/"add tests later". diff --git a/docs/superpowers/specs/2026-08-26-umbrella-backport-8060s-design.md b/docs/superpowers/specs/2026-08-26-umbrella-backport-8060s-design.md new file mode 100644 index 0000000..af5440f --- /dev/null +++ b/docs/superpowers/specs/2026-08-26-umbrella-backport-8060s-design.md @@ -0,0 +1,179 @@ +# Umbrella repo + 8060S backport — Design + +Date: 2026-08-26 +Status: implemented — wave-1 shipped as PR #1 (executed 2026-08-26; per-task reviews + final whole-branch review clean) + +## Goal + +Merge the valuable contributions of +[`SenseNova-U1.5-ROCm-8060S`](https://github.com/AIwork4me/SenseNova-U1.5-ROCm-8060S) +(the Strix Halo / gfx1151 sibling repo) into this repo, and establish a +durable two-repo collaboration model, **while keeping the 8060S repo +independent** — its own identity, quickstart UX, evidence corpus, and +release cadence. + +This repo becomes the **multi-hardware umbrella**: community-facing, +hardware-generic assets plus per-hardware profiles. The 8060S repo becomes +the **Strix Halo lab**: first to receive APU-specific work; verified +experiments graduate upstream into this repo via PR. + +## Research findings that shaped the design + +1. **The two repos are same-origin siblings, not forks.** Both were created + 2026-08-22 (one hour apart), never git-linked, and carry the same + `patches/0001` (interleave image_size hotfix — ours regenerated to the + upstream PR#260 head e2f2c865, theirs still the 76c32c2-era format-patch). + Git-level unification (forking / history grafting) is therefore both + painful and pointless: the trees diverged structurally from day one. +2. **Structural divergence map.** + + | | this repo (gfx1100) | 8060S (gfx1151) | + |---|---|---| + | default branch | `main` | `master` | + | python code | `src/sensenova_u1_rocm/` | `python/` (loose tools) | + | scripts | numbered pipeline + `scripts/lib/common.sh` | `run-*.sh` per task + root `quickstart.sh` | + | evidence | `docs/results/` receipts (receipt.py, sha256) | `evidence/` (results.json per run, 178 files) | + | quality gates | CI (shellcheck strict, offline pytest), CONTRIBUTING, CHANGELOG, tests | none of these | + +3. **What 8060S has that we lack** (verified by reading its tree): + a hardware-independent `torch.compile`/cudagraph-safety patch + (env-gated, V1–V8 verified), `gpu_monitor.py` (VRAM **and GTT** sampling — + we sample neither), `rocm_check.py` (env sanity incl. APU GTT-spill + alloc test), `ROCm_PORTING.md` (mixed generic + APU knowledge), + `hsa_fix.sh` (Strix Halo wheel-HSA crash workaround), and the AOTriton + hd72 root-cause chain (their issue #54, mirrors our pytorch#194498 work). +4. **Overlap to avoid**: their `bench_report.py` duplicates our + `summarize_results.py` purpose (aggregate run results into a summary + table) — porting both would create a two-headed tool. +5. **Both repos pin the same upstream checkout** (`feat/u1.5` @ `76c32c2`, + patches applied by the venv setup script) — so their patch applies (or + regenerates cleanly) against our pin with high probability; must be + verified, not assumed. +6. **The working machine for this effort is the 8060S reference host** + (Ryzen AI MAX+ PRO 395, gfx1151). gfx1100 re-validation is possible only + later, on that host. + +## Decisions (user-approved 2026-08-26) + +| decision point | choice | +|---|---| +| merge scope | umbrella repo: generic assets in-tree + Strix Halo as a hardware profile; 8060S keeps evidence/announcements/showcase | +| structure relation | 8060S aligns to this repo's structure over time (its own pace, three steps) | +| evidence policy | layered attribution: merge with `verified on: gfx1151` labels linking 8060S evidence; drop labels only after gfx1100 receipts exist | +| execution scope | design + wave-1 executed locally (branch, port, test on gfx1151); push/PR only after user confirmation; gfx1100 re-validation deferred | + +## Design + +### 1. Governance and repo relationship + +- This repo: **no rename** (already hardware-generic), description and + README scope broadened to "AMD GPUs" covering gfx1100 dGPU + gfx1151 APU. + New `docs/hardware/strix-halo/` profile directory; existing gfx1100 + content stays where it is (minimal churn). +- 8060S keeps: independent repo identity, root `quickstart.sh` one-command + UX, its `evidence/` corpus, announcements, showcase, own releases. +- Cross-linking: 8060S README gains an explicit "upstream: this repo" note; + this repo's README hardware table links both. GitHub topics unified + (`rocm`, `amd`, `sensenova`, …). +- Relationship model: classic downstream–upstream. 8060S experiments + graduate into this repo via PR when they meet the graduation criteria + (§5); this repo's shared-dir changes flow down to 8060S via cherry-pick. + +### 2. Wave-1 backport inventory + +Ordered by value × independence; every item lands as its own commit with +`Ported from AIwork4me/SenseNova-U1.5-ROCm-8060S@`. + +| # | from 8060S | to this repo | action | +|---|---|---|---| +| 1 | `patches/0002-optional-torch-compile…patch` | `patches/0003-…` (**renumbered**; our 0002 is SDPA) | adapt to our patch format; verify clean apply on pin `76c32c2`; smoke `SENSENOVA_COMPILE=1` on gfx1151 | +| 2 | `python/gpu_monitor.py` | `scripts/gpu_monitor.py` | adopt script conventions (stdlib-only, usage(), `set -euo pipefail` n/a for py); add offline unit tests; GTT sampling is new capability | +| 3 | `python/rocm_check.py` | `scripts/rocm_check.py` | generic env-check logic; alloc-test section labelled APU/GTT-scoped | +| 4 | `docs/ROCm_PORTING.md` generic sections | `docs/porting.md` (standalone new file — the source has 7 substantive sections; folding into getting-started would bloat it) | split: generic → here, APU-specific → `docs/hardware/strix-halo/` | +| 5 | `scripts/hsa_fix.sh` + GTT-spill knowledge | `docs/hardware/strix-halo/` (doc + script) | hardware-profile absorption | +| 6 | AOTriton #54 / V2 root-cause chain | `docs/upstream/` entry | alongside pytorch#194498 tracking | +| 7 | `python/bench_report.py` | — | **deferred** (overlaps `summarize_results.py`) | +| 8 | `evidence/`, announcements, showcase | — | **not merged**; referenced by link | + +Docs ported in wave 1 ship EN + CN variants together (bilingual discipline). + +Wave 1 also includes two **this-repo-side adaptations** (not ports, listed +here so the plan is exhaustive): the `receipt.py` `hardware` field with its +offline test (§3), and the governance additions (CoC, SECURITY.md, issue/PR +templates, CONTRIBUTING cross-repo section, README hardware-table update). + +### 3. Evidence and attribution rules + +- Layered labels: merged hardware-independent claims carry + `verified on: gfx1151 (Strix Halo) — evidence: <8060S link>` until a + gfx1100 receipt lands in `docs/results/`, then the label is removed. +- `receipt.py` gains an optional `hardware` field (e.g. `"gfx1151"`) so + cross-hardware evidence is machine-queryable; offline unit test included. +- Attribution: every ported commit carries `Ported from …@`; original + authors get `Co-authored-by:` where applicable. 8060S records + `Merged-upstream: ` on its side — bidirectional trace. + +### 4. 8060S alignment roadmap (executed in the 8060S repo, its own pace) + +1. **Governance**: CONTRIBUTING (incl. graduation flow), CHANGELOG + (backfilled from git history), Code of Conduct, SECURITY.md, issue/PR + templates — adapted from this repo, bilingual. +2. **Structure**: `python/*` → `scripts/*` with this repo's names (after + wave 1 lands here, 8060S adopts the merged versions verbatim — sync cost + drops to cherry-pick); root `quickstart.sh` remains as a thin wrapper. +3. **Quality gates**: this repo's CI (shellcheck strict + offline CPU + pytest) and a `tests/` skeleton. + +Also in this repo (governance gaps found in self-audit): Code of Conduct, +SECURITY.md, issue/PR templates, GitHub Discussions enabled (8060S already +uses Discussions), CONTRIBUTING gains a cross-repo contribution section. + +### 5. Ongoing sync process (bidirectional) + +- **8060S → here (graduation)**: experiment verified on gfx1151 → meets + graduation criteria → PR to this repo labelled `backport`, CI green → + merged → 8060S records `Merged-upstream`. +- **Here → 8060S (flow-down)**: changes to shared dirs (`patches/`, + `scripts/`) are cherry-picked / replayed by 8060S as needed. +- **Graduation criteria (written down)**: ① hardware-independent or + absorbed as a hardware profile ② gfx1151 evidence exists ③ passes this + repo's offline CI ④ EN+CN docs complete ⑤ no two-headed tool conflict. +- Each CONTRIBUTING documents the other direction; PRs/issues cross-linked + between repos. + +### 6. Testing, risks, execution boundary + +Wave-1 tests (all doable on the gfx1151 host): +- `git apply --check` of the renumbered torch.compile patch against the + pinned upstream checkout; real `SENSENOVA_COMPILE=1` smoke run. +- `gpu_monitor.py` / `rocm_check.py` real runs on gfx1151 + new offline + tests under `tests/`. +- This repo's `python3 -m pytest tests/ -q` green + shellcheck clean + (same gates as CI). + +Risks & mitigations: +- torch.compile patch drift vs our pin → regenerate first, verify apply, + rework in our format if conflicting. +- two-headed tools → bench_report deferred. +- doc-reorg churn → zero moves of existing content, only new dirs. +- network: git transport to github.com currently failing from this host + (TLS resets); work proceeds on local clones initialised from `gh api` + tarballs; push/PR deferred until both network and user confirmation. + +## Non-goals + +- No repo rename, no fork conversion, no git-history grafting. +- No bulk import of 8060S `evidence/` (178 files) or announcements. +- No 8060S restructuring *inside* this repo's PRs — that happens in the + 8060S repo on its own schedule. +- No gfx1100 re-validation inside wave 1 (deferred; needs that host). + +## Success criteria + +1. Wave-1 items 1–6 merged as individually-attributed commits; CI green; + bilingual docs in place. +2. Governance gaps closed in this repo (CoC, SECURITY, templates, + Discussions, CONTRIBUTING cross-repo section). +3. Graduation criteria + sync process written into both CONTRIBUTINGs. +4. 8060S still fully functional and independent (quickstart UX unchanged), + with its alignment roadmap committed as its own docs. diff --git a/docs/upstream/aotriton-54.md b/docs/upstream/aotriton-54.md new file mode 100644 index 0000000..041b8b8 --- /dev/null +++ b/docs/upstream/aotriton-54.md @@ -0,0 +1,101 @@ +# Upstream: AOTriton flash SDPA silently wrong for head_dim=72, seq>1024, non-causal (ROCm/aotriton#54 — fix unmerged) + +verified on: gfx1151 (Strix Halo / Radeon 8060S); no gfx1100 receipt yet. +The defect is kernel-level, so any host whose torch wheel bundles the +unfixed AOTriton release line is presumed affected until proven otherwise. + +## Summary + +The AOTriton flash SDPA forward kernel returns **silently wrong output** +for one narrow but real configuration class — no launch error, no +exception, just garbage (worse than a crash, because nothing surfaces): + +- head_dim = **72** exactly (64/80/88/96/104/112/120/128 all measured OK) +- seq_len **> 1024** (1024 OK; 1025 already fails) +- **non-causal** attention only (causal at the same shape is correct — + a different tile configuration) + +This is the shape of typical ViT patch attention: the Qwen-Image-Bench +judge ViT is 16 heads × hidden 1152 (= head_dim 72) and a 1024² image +yields 4096 patches — squarely inside the envelope. SenseNova's own +attention shapes (head_dim 64/128) are **unaffected**; that is why +AOTriton-enabled SenseNova runs stay correct while the judge ViT sharing +the same venv produced garbage ("乱码") — the observation that opened this +investigation. + +Note the contrast with the gfx1100 fused-backend *launch* failure tracked +in [issue-pytorch-rocm-sdpa-backends.md](issue-pytorch-rocm-sdpa-backends.md): +this kernel launches fine and lies. + +## Symptoms (speedup2 re-verification, 2026-08-25/26) + +- relative error vs the MATH backend (bf16, b=1/16 heads): + **1.6e-01 … 7.3e+00 (≈16–730%)** — dominantly *finite* garbage, not NaN +- error magnitude **drifts run-to-run even at fixed seed** (single-run + snapshots are not an invariant; one re-verification capped hd72@4096 at + 4.9e-01) +- **sporadic NaN in ~1/5 of calls** (5 consecutive same-input calls: + [3.9e-1, nan, 5.5e-1, 1.4e-1, 3.4e-1]) — "not NaN" is not an invariant; + report as interval + occasional NaN +- the **LSE itself is corrupted** (softmax denominator polluted); errors + spread across the whole sequence +- reproduces via direct `aten::_scaled_dot_product_flash_attention` — + kernel-level, not an SDPA dispatch artifact +- fp16 also affected; bad across the whole batch × heads matrix + (occupancy-independent — split-KV hypothesis excluded) + +## Root cause and upstream status + +Upstream [ROCm/aotriton issue #54](https://github.com/ROCm/aotriton/issues/54) +(open since 2024-11): the fwd kernel's +`qk += (Qk_scale * tl.dot(q0, k0))` lacks an explicit +`out_dtype=tl.float32`, so on gfx1151 the dot result overflows / drops +mantissa in bf16/fp16 before the scale multiply → garbage or NaN. +head_dim 72 is the size that takes the BLOCK_DMODEL tiled-dot path +(the `BLOCK_DMODEL1>0` branch) — the necessary condition to reach the +defective line. + +- Fix commit `8232d69672` (2026-06-01, branch + `fix/gfx1151-bf16-dot-fp32-accumulator`) — **still unmerged**: main and + every release tag 0.11.2b…0.13.50tp carry the unfixed code; the tested + wheel bundles `libaotriton_v2.so.0.11.2` (unfixed). The fix's commit + notes cite the BLOCK_M=128/BLOCK_N=64 trigger, matching the measured + envelope. +- Production repro + failure envelope + causal control reported upstream + on 2026-08-26 + ([issuecomment-5420432049](https://github.com/ROCm/aotriton/issues/54#issuecomment-5420432049)), + requesting that the fix branch be merged and released. This is an + upstream defect — nothing to patch in this repo; this entry tracks it + until a release carries `8232d69672`. + +## Cross-stack verification (gfx1151) + +| stack | hd=72 result | +|---|---| +| TheRock `torch 2.10.0+rocm7.13.0a20260513` (HIP 7.13.26183, vLLM venv) | seq ≥1088 BAD (1.6e-01…7.3e+00); hd 64/128 and hd72@1024 OK | +| AMD official `torch 2.12.0+rocm7.14.0` (HIP 7.14.60850) | seq 1088/2048/4096 → rel err 1.9e-1 / 3.2e-1 / 9.1e+0, all BAD; hd 64/128 and hd72@1024 OK | + +The defect spans at least ROCm 7.13-alpha (TheRock) through 7.14.0 +(official). (The V2 note initially mis-attributed the repro to +2.12+rocm7.14; corrected 2026-08-26 — attribution above is final.) + +## Workaround until fixed upstream + +SenseNova's own attention (head_dim 64/128) needs no action. For +**foreign models with other head dims (hd≠64/128 — classically a +16-head/1152-hidden ViT)** run in a venv where AOTriton is enabled: + +- **disable AOTriton for that model / pin SDPA to the MATH backend** + (the 8060S repo's judge service keeps AOTriton off as its standing + policy; the run scripts there disable it with `SENSENOVA_NO_AOTRITON=1`) +- do not treat "no NaN" as evidence of health — finite garbage dominates + +## Evidence + +- Verification chain (V2 root-cause closure, failure envelope, upstream + filing; commits `d0af6efc` / `987015a1` / `28e4f3af`): + [evidence/speedup2/VERIFICATION.md §V2](https://github.com/AIwork4me/SenseNova-U1.5-ROCm-8060S/blob/master/evidence/speedup2/VERIFICATION.md) + (directory: + [evidence/speedup2/](https://github.com/AIwork4me/SenseNova-U1.5-ROCm-8060S/tree/master/evidence/speedup2)) +- User-facing caveat (exact guidance quoted above): + [docs/TROUBLESHOOTING.md §5](https://github.com/AIwork4me/SenseNova-U1.5-ROCm-8060S/blob/master/docs/TROUBLESHOOTING.md) diff --git a/patches/0003-optional-torch-compile-and-cudagraph-safety.patch b/patches/0003-optional-torch-compile-and-cudagraph-safety.patch new file mode 100644 index 0000000..741c194 --- /dev/null +++ b/patches/0003-optional-torch-compile-and-cudagraph-safety.patch @@ -0,0 +1,60 @@ +diff --git a/examples/t2i/inference.py b/examples/t2i/inference.py +index 63a7416..20a5e95 100644 +--- a/examples/t2i/inference.py ++++ b/examples/t2i/inference.py +@@ -2,11 +2,13 @@ from __future__ import annotations + + import argparse + import json ++import os + from pathlib import Path + from typing import Sequence + + import numpy as np + import torch ++import torch._dynamo + from PIL import Image + + import sensenova_u1 +@@ -428,6 +430,19 @@ def main() -> None: + print(f"load lora {args.lora_path}") + engine.model = load_and_merge_lora_weight_from_safetensors(engine.model, args.lora_path) + ++ comp_mode = os.environ.get("SENSENOVA_COMPILE", "") ++ if comp_mode: ++ if os.environ.get("SENSENOVA_COMPILE_SUPPRESS", "0") == "1": ++ torch._dynamo.config.suppress_errors = True ++ comp_kwargs: dict = {"dynamic": False} ++ if comp_mode not in ("1", "default"): ++ comp_kwargs["mode"] = comp_mode ++ engine.model._t2i_predict_v = torch.compile(engine.model._t2i_predict_v, **comp_kwargs) ++ if os.environ.get("SENSENOVA_COMPILE_FEATURES", "0") == "1": ++ engine.model.extract_feature = torch.compile(engine.model.extract_feature, **comp_kwargs) ++ engine.model.patchify = torch.compile(engine.model.patchify, **comp_kwargs) ++ print(f"[compile] enabled {comp_kwargs}") ++ + cfg_interval = tuple(args.cfg_interval) + + if args.prompt is not None: +diff --git a/src/sensenova_u1/models/neo_unify/modeling_neo_chat.py b/src/sensenova_u1/models/neo_unify/modeling_neo_chat.py +index f52a7fe..47bda48 100644 +--- a/src/sensenova_u1/models/neo_unify/modeling_neo_chat.py ++++ b/src/sensenova_u1/models/neo_unify/modeling_neo_chat.py +@@ -1815,9 +1815,16 @@ class NEOChatModel(PreTrainedModel): + timestep_embeddings += noise_embeddings + image_embeds = image_embeds + timestep_embeddings + ++ # no-op without cudagraphs; prevents graph-output overwriting when ++ # the same compiled graph runs twice per step (cond + uncond) ++ torch.compiler.cudagraph_mark_step_begin() + v_pred_condition = self._t2i_predict_v(image_embeds, indexes_image_condition, attention_mask_condition, past_key_values_condition, t, z, image_token_num=token_h*token_w, timestep_embeddings=timestep_embeddings, image_size=image_size) +- ++ # materialize outside the cudagraph output pool before the same ++ # compiled callable re-runs for the CFG unconditional pass ++ v_pred_condition = v_pred_condition.clone() ++ + if t >= cfg_interval[0] and t <= cfg_interval[1] and cfg_scale > 1: ++ torch.compiler.cudagraph_mark_step_begin() + v_pred_uncondition = self._t2i_predict_v(image_embeds, indexes_image_uncondition, attention_mask_uncondition, past_key_values_uncondition, t, z, image_token_num=token_h*token_w, timestep_embeddings=timestep_embeddings, image_size=image_size) + if cfg_norm == 'cfg_zero_star': + positive_flat = v_pred_condition.view(batch_size, -1) diff --git a/patches/README.md b/patches/README.md index 58bb877..c6000a3 100644 --- a/patches/README.md +++ b/patches/README.md @@ -78,3 +78,21 @@ on torch 2.8+BLAS-preload (receipt `docs/results/validation/t2i-torch212-fixed.j still open; [pytorch/pytorch#194498](https://github.com/pytorch/pytorch/issues/194498); compatibility write-up: [OpenSenseNova/SenseNova-U1#261](https://github.com/OpenSenseNova/SenseNova-U1/issues/261)). + +## 0003-optional-torch-compile-and-cudagraph-safety.patch + +**Opt-in acceleration (all platforms).** Adds env-gated `torch.compile` +around `_t2i_predict_v` (and optionally `extract_feature`/`patchify`): +`SENSENOVA_COMPILE=1|default|reduce-overhead|…` enables, +`SENSENOVA_COMPILE_SUPPRESS=1` sets dynamo suppress-errors, +`SENSENOVA_COMPILE_FEATURES=1` also compiles the feature/patchify path. +The model-side hunk marks cudagraph step boundaries and materializes the +conditioned prediction before the unconditioned re-run of the same +compiled callable — without it, cudagraph output pooling overwrites the +CFG condition with the uncondition result. + +verified on: gfx1151 (Strix Halo) — V1–V8 speedup + Qwen-Image-Bench +paired quality validation live in the +[8060S repo](https://github.com/AIwork4me/SenseNova-U1.5-ROCm-8060S) +(`evidence/speedup2/`, PASS-with-caveat: compile stays opt-in until a +gfx1100 receipt exists here). diff --git a/scripts/gpu_monitor.py b/scripts/gpu_monitor.py new file mode 100755 index 0000000..0a5510a --- /dev/null +++ b/scripts/gpu_monitor.py @@ -0,0 +1,100 @@ +#!/usr/bin/env python3 +"""Sample AMD GPU memory usage (VRAM + GTT) while another process runs. + +Usage: + python scripts/gpu_monitor.py --interval 0.5 --out docs/results/logs/mem_usage.jsonl + +Writes one JSON object per sample: + {"ts": 1724..., "vram_used_gib": 12.3, "gtt_used_gib": 4.1} + +Run it in the background, start the workload, then send SIGINT/SIGTERM (or +call stop()) — it prints a summary (peak VRAM / peak GTT / peak total) and +exits. The summary line is also appended to the output file for easy +aggregation by scripts/validate.sh receipts. +""" + +from __future__ import annotations + +import argparse +import json +import signal +import subprocess +import sys +import time +from pathlib import Path + + +def _used_bytes(which: str) -> int | None: + """rocm-smi --csv emits: device,