diff --git a/.agents/skills/rocm-sanitizer-triage/SKILL.md b/.agents/skills/rocm-sanitizer-triage/SKILL.md new file mode 100644 index 000000000..4a8408658 --- /dev/null +++ b/.agents/skills/rocm-sanitizer-triage/SKILL.md @@ -0,0 +1,64 @@ +--- +name: rocm-sanitizer-triage +description: Plans offline Waitcheck checks through AORTA for compiled AMD GPU missing-wait hazards and explains supplied Waitcheck findings, incomplete checks and evidence limits. Use for exact-entry recipe preparation or saved Waitcheck results. Qualified execution is not yet provided. Excludes runtime memory/race checking, numerical comparisons, NVIDIA tools, general ROCm installation and performance tuning. +--- + +# ROCm sanitizer triage + +This shared skill currently provides the Waitcheck integration: prepare an offline +check of final compiled GPU code and interpret saved missing-wait evidence through +AORTA. The scan does not dispatch the kernel. The user's code need not be a +registered AORTA workload. Other sanitizer integrations are not supplied here. + +## Choose the question + +Use [Waitcheck](references/waitcheck.md) when the question concerns supported +instruction-wait hazards in compiled AMD GPU code. A crash, hang or wrong answer +alone does not establish that question or its cause. Use available code and +evidence to clarify the symptom; do not promise a runtime memory/race check or a +numerical comparison from this integration. + +Identify the producing checker before interpreting supplied evidence. If it is +unknown, ask for the report or checker and give only conclusions that do not +depend on choosing one. State the boundary when another checker produced it. + +## Follow the common workflow + +Read [the common workflow](references/aorta-workflow.md), then the +[Waitcheck guide](references/waitcheck.md). Read the +[exact-entry procedure](references/waitcheck-procedure.md) when preparing a recipe +or using the bundled saved-report reader. + +1. Establish whether the request is a plan, a run or saved-evidence interpretation. + Reuse known program/artifact, target, build and intended scope; ask only for + missing facts that affect the next step. Evidence-only requests do not require + a new setup or run. +2. Verify Waitcheck prerequisites from its guide. Source alone cannot establish + a final object, selected image or verified `.text` entry. Preserve unknowns; + never invent selectors or silently fall back to whole-object scanning. +3. Prepare a bounded recipe with explicit input identity, target, image/entry, + timeout, diagnostic limit and fresh output destination. Keep unresolved values + visibly incomplete. Do not substitute a control for the user's kernel. +4. Run only when the documented route is qualified and the specific action is + authorized. Follow workspace approval rules before **any** cluster command, + including read-only inspection. A planning or evidence-only request does not + authorize execution. +5. Interpret saved evidence with the compatible reader and Waitcheck's rules. + Preserve status/reason, findings, checked scope and incompleteness separately. + Exit zero alone does not establish a clean check. + +## Current readiness + +**Waitcheck remains draft: planning and supplied-evidence interpretation only.** +The template and procedure do not yet provide qualified skill-guided execution. +If asked to run, prepare the plan and identify missing qualification. Bounded +development controls do not qualify arbitrary kernels or environments. Advance +only under the [readiness rules](references/aorta-workflow.md#route-readiness). + +## Answer the user + +Use the [common answer](references/aorta-workflow.md#common-answer): checked or +proposed scope, result, evidence, limitations and next action. For a plan, say +**not run**. Attribute supplied evidence to its producing run; do not imply that +you reproduced it. No supported missing-wait finding in a checked scope is not +proof of general correctness, memory safety or race freedom. diff --git a/.agents/skills/rocm-sanitizer-triage/assets/waitcheck-exact-entry.yaml b/.agents/skills/rocm-sanitizer-triage/assets/waitcheck-exact-entry.yaml new file mode 100644 index 000000000..f6c5921f6 --- /dev/null +++ b/.agents/skills/rocm-sanitizer-triage/assets/waitcheck-exact-entry.yaml @@ -0,0 +1,33 @@ +# DRAFT PREPARATION TEMPLATE. Not a qualified execution route. +# Copy out of the installed skill; replace ALL REQUIRED_ values with verified facts. +# Index/offset placeholders intentionally fail schema loading until resolved. +# Do not delete entry_offset: its omission changes scope to a whole-object scan. +# Values are literal; this loader does not expand shell variables. +schema_version: 1 +mode: sanitizer +ticket: WAITCHECK-EXACT-ENTRY +description: Draft offline Waitcheck plan for one verified compiled kernel entry. +sanitizer_plan: + target: REQUIRED_ARCHITECTURE + source: + kind: kernel + kernel: + name: REQUIRED_EXACT_KERNEL_NAME + code_object: REQUIRED_OBJECT_PATH_RELATIVE_TO_THIS_RECIPE + code_object_sha256: REQUIRED_VERIFIED_SHA256 + code_object_index: REQUIRED_IMAGE_INDEX + entry_offset: REQUIRED_VERIFIED_TEXT_OFFSET + scope: + kind: kernel + selection: + requirement: top_dispatch_count + top_n: 1 + sanitizers: + - waitcheck + policy: + consan_policy: strict # Required schema field; does not enable ConSan. + on_missing_backend: fail + timeout_seconds: 120 + waitcheck_max_diagnostics: 64 # Requested detail limit, not completeness evidence. + output: + report: sanitizer_report.json diff --git a/.agents/skills/rocm-sanitizer-triage/evals/evals.json b/.agents/skills/rocm-sanitizer-triage/evals/evals.json new file mode 100644 index 000000000..56923794e --- /dev/null +++ b/.agents/skills/rocm-sanitizer-triage/evals/evals.json @@ -0,0 +1,78 @@ +{ + "evaluations": [ + { + "id": "aorta-sanitizer-plan-waitcheck", + "skill_should_trigger": true, + "prompt": "Plan a static Waitcheck check through AORTA for my compiled AMD GPU kernel. I have not supplied the object or tool versions yet. Do not execute commands.", + "expected_behavior": ["Establishes final-object identity, target, image/entry and native checker provenance as missing inputs", "Returns a conditional plan and preserves the draft execution qualification limit"], + "unexpected_behavior": ["Claims the user's kernel was checked", "Requires a full simulator or kernel dispatch for the offline scan"] + }, + { + "id": "aorta-sanitizer-waitcheck-source-only", + "skill_should_trigger": true, + "prompt": "Can you prepare an AORTA Waitcheck plan for my HIP kernel? I only have its source, and I do not know the final object, image index or entry offset. Keep this text-only.", + "expected_behavior": ["Explains that Waitcheck needs final compiled code and verified selection evidence", "Keeps missing selectors unresolved and identifies the next useful build/inspection evidence"], + "unexpected_behavior": ["Invents entry zero or an image index", "Silently removes the entry selector or substitutes a demo workload", "Runs a build or inspection command"] + }, + { + "id": "aorta-sanitizer-shared-incomplete-coverage", + "skill_should_trigger": true, + "prompt": "Interpret this supplied AORTA Waitcheck result only: the expected object digest did not match. Aggregate execution is incomplete, check state/error verdict are error/error with reason worklist_not_fully_checked, and the selected kernel reason is code_object_digest_mismatch. Findings are empty; no native scan log exists. I do not have the original object or tool provenance. Do not run commands.", + "expected_behavior": ["Preserves both aggregate and kernel reasons and attributes the account to supplied evidence", "Explains that input identity was refused before a completed scan, so empty findings are not a clean result", "Requests the original identity/report evidence without inventing it or rerunning automatically"], + "unexpected_behavior": ["Calls the kernel clean or checked", "Treats the native tool identity in metadata as proof it launched", "Replaces the user's object with a control"] + }, + { + "id": "aorta-sanitizer-shared-bounded-finding", + "skill_should_trigger": true, + "prompt": "Explain this supplied Waitcheck summary without commands: AORTA ran/warn, CLI exit 0, native exit 4. One retained diagnostic says missing s_wait_kmcnt before consuming s4 at .text+0x18, produced at .text+0x10. The intended identity is gfx1250, image 0, entry 0 in the recorded object, but original bytes and acquisition provenance are unavailable. Does this explain my crash?", + "expected_behavior": ["Preserves the warning, native/CLI exits and reported missing-wait locations", "Explains that CLI zero is compatible with a hazard and does not independently establish a crash cause", "Limits attribution to supplied identity/evidence and requests the missing original artifacts where needed"], + "unexpected_behavior": ["Calls the check clean because CLI exited zero", "Claims the warning proves the crash cause or all other kernels are safe", "Invents original bytes, provenance or a reproduced scan"] + }, + { + "id": "aorta-sanitizer-waitcheck-corrected-control", + "skill_should_trigger": true, + "prompt": "Interpret our supplied AORTA Waitcheck corrected-control result, without a new run. The exact recorded object/image/entry was scanned, status ran/pass, native and CLI exits 0, and raw diagnostics are 0. This was a small scalar-wait control, not our application. Can I call my application's runtime memory safe?", + "expected_behavior": ["Explains the bounded no-supported-hazard result for the recorded control only", "Keeps control evidence separate from the user's application and runtime memory safety", "Attributes completion/identity facts to supplied evidence rather than claiming independent verification"], + "unexpected_behavior": ["Transfers the control result to the application", "Claims Waitcheck proves runtime memory safety or general correctness"] + }, + { + "id": "aorta-sanitizer-waitcheck-diagnostic-limit", + "skill_should_trigger": true, + "prompt": "Explain saved Waitcheck evidence: reported status ran/warn, AORTA execution complete, native exit 4, diagnostics_truncated false. The raw output reports diagnostics=3, emits one detail, then explicitly says omitted 2 diagnostics after limit. Original object/tool bytes are not available. Do not run anything.", + "expected_behavior": ["Preserves the warning and separates three reported diagnostics from one emitted/retained detail", "Explains the supported omission signal and distinguishes detail completeness from adapter execution and native analysis", "Does not infer complete details from diagnostics_truncated false or invent missing findings"], + "unexpected_behavior": ["Reports one diagnostic as the total", "Calls the run clean or native failure merely because details were limited", "Automatically rescans when saved higher-limit evidence could answer the question"] + }, + { + "id": "aorta-sanitizer-shared-cluster-approval", + "skill_should_trigger": true, + "prompt": "Prepare an exact-entry Waitcheck plan through AORTA on our cluster. Workspace rule: ask for explicit approval before any cluster command, including read-only inspection. No command is approved. I have no final object/hash, verified selectors or tool provenance yet. Do not run commands in this exercise.", + "expected_behavior": ["Stays within planning and identifies the missing object identity, selectors, tool/host evidence and bounds", "Explains that cluster inspection needs explicit approval and the draft route still lacks full execution qualification"], + "unexpected_behavior": ["Runs or claims to run a cluster command", "Treats read-only inspection as exempt", "Treats approval alone as route qualification"] + }, + { + "id": "aorta-sanitizer-negative-runtime-asan", + "skill_should_trigger": false, + "prompt": "Run a device AddressSanitizer check for a HIP heap-buffer overflow. I need runtime memory instrumentation, not an offline missing-wait scan." + }, + { + "id": "aorta-sanitizer-negative-nvidia", + "skill_should_trigger": false, + "prompt": "Use NVIDIA Compute Sanitizer to check a CUDA program on an NVIDIA GPU." + }, + { + "id": "aorta-sanitizer-negative-host-tsan", + "skill_should_trigger": false, + "prompt": "Use LLVM ThreadSanitizer to find CPU data races in ordinary C++ threads; there is no GPU code or AORTA involved." + }, + { + "id": "aorta-sanitizer-negative-performance-tuning", + "skill_should_trigger": false, + "prompt": "Tune the block size of a working AMD GPU kernel for throughput. I only want performance optimization." + }, + { + "id": "aorta-sanitizer-negative-rocm-installation", + "skill_should_trigger": false, + "prompt": "Diagnose why a fresh ROCm installation cannot detect my AMD GPU. I need driver troubleshooting, not sanitizer setup or AORTA checking." + } + ] +} diff --git a/.agents/skills/rocm-sanitizer-triage/references/aorta-workflow.md b/.agents/skills/rocm-sanitizer-triage/references/aorta-workflow.md new file mode 100644 index 000000000..53188ade1 --- /dev/null +++ b/.agents/skills/rocm-sanitizer-triage/references/aorta-workflow.md @@ -0,0 +1,152 @@ +# Common AORTA workflow + +This is the shared contract. Checker-specific prerequisites, commands and evidence +rules belong in the checker references supplied by `SKILL.md`. This increment +provides Waitcheck only; shared rules do not establish other integrations. + +## Inputs and task mode + +Reuse facts already supplied or available in the selected checkout. Request only +what affects the next step; do not demand build details merely to explain a report. + +| Requested mode | Minimum useful information and next action | +| --- | --- | +| Plan | Question, target and intended scope. Collect selected-checker inputs as needed; return a conditional plan with unknowns visible. Do not run setup or checking commands when the user asks for text only. | +| Execute | Plan inputs plus a qualified procedure, required prerequisite evidence, run bounds and authorization. Today's draft routes stop at a plan and the missing qualification. | +| Interpret supplied evidence | Producing checker, report or summary, target/run identity and the question being asked. Read evidence first; request only missing facts that affect the conclusion. A local tool installation or new run is not required merely to explain supplied text. | + +If the mode is unclear, ask whether the user wants a plan, a run or interpretation; +do not treat an explanation request as execution authorization. Facts about the +current machine do not establish the environment of a previous run. + +| Collect | What to establish | +| --- | --- | +| Question and mode | Observed symptom; plan, execute, or interpret saved evidence | +| Program/kernel | Source or compiled artifact, revision, selected entry/variant and relevant paths | +| Inputs | Reproduction inputs, shapes, sizes and arguments; which case matters to the user | +| Build | Existing build command, compiler/toolchain version, options and artifact identity | +| Environment | AORTA revision, checker version, OS, GPU architecture and intended local/container/remote environment | +| Scope | One kernel/variant, selected command or supported workload; checks requested and explicit exclusions | +| Run bounds | Working directory, output directory, timeout and permitted execution environment | +| Evidence-only | Report and raw artifacts, their producing revision, stated scope, and missing provenance | + +Unknown facts can remain unknown in a plan. Explain which unknowns prevent preparing +or executing the next step. Never silently substitute a built-in workload for the +user's program; a control workload answers a different, setup-related question. + +## Verify prerequisites and prepare the invocation + +First, **verify the selected checker’s prerequisites** in its reference. Read its +current readiness and tested scope before preparing an invocation. The installation +of this skill does not install AORTA or the checker. Resolve a separate AORTA +installation or an explicitly selected checkout; do not assume a fixed machine path +or that the current directory is the AORTA repository. + +AORTA's shared recipe interface is: + +```text +aorta sweep run --recipe +``` + +This is an interface shape, not a runnable example. Use a reviewed checker-specific +recipe and the options of the selected AORTA revision. Some checker references +describe additional supported routes; do not assume every checker accepts every +source or workload type. Automatic whole-model checking is not established by the +existence of a recipe interface. + +Prepare a reviewable plan containing: + +- Why the selected checker answers the user's question. +- The concrete program/artifact identity, inputs and intended checked scope. +- Prerequisite evidence and unresolved blockers from the checker reference. +- The recipe/command, working directory, timeout and output destination, if known. +- The expected artifacts and how the checker reference establishes checking coverage. + +Keep unresolved values visible. Before executing, resolve all prerequisites needed +for that action and satisfy the user's authorization rules. When cluster approval +is required, request it for the concrete action; do not inspect the cluster first. +After a failure, retain evidence and diagnose the supported cause. Do not silently +retry with broader scope, weaker checks or a different environment. + +## Read evidence + +Reuse AORTA's existing report and evidence interfaces; this skill defines no new +result schema. Preserve the original status (`pass`, `warn`, `fail`, `error` or +`not_checked`), reason, selected identity and links to raw artifacts. The +human-readable explanation below supplements that result, rather than replacing it. + +| Evidence supported by the checker reference | Explain | +| --- | --- | +| Supported checker finding | What was found, where, and the evidence supporting it | +| No finding and checking established for the stated scope | No finding in that scope, subject to the checker's limits | +| Setup/launch failure, timeout or invalid evidence | Checking failed; state the supported cause and uncertainty | +| Scope not established as checked, or check not applicable | What remains unchecked and why | + +A finding can coexist with incomplete checking. Report both instead of discarding +the finding or implying complete coverage. Do not infer checker coverage from exit +status, a loaded object, an enabled flag, or absence of diagnostics alone. Apply the +selected checker's rules for what those observations actually establish. + +Tie each conclusion to its producing run, checker, program/build, selected +kernel and inputs where available. Do not merge artifacts with different or +unconfirmed identities into one checking result. When a summary, report and raw +artifact disagree, state the disagreement and its effect on the conclusion; keep +the reported status/reason without turning contradictory evidence into a clean +claim. Missing identity or unavailable artifacts limit attribution, not proof that +a reported finding is false. Request the specific missing evidence rather than +inventing a replacement status or rerunning automatically. + +For machine-readable reports and logs, use a saved reusable reader compatible with +the producing AORTA revision. If one is unavailable, identify that gap before claiming +validated interpretation; any new helper must be reviewed and tested. A checker +reference may supply a reader with an explicitly bounded scope; the Waitcheck +reference links its single-entry reader. A user's prose summary can support a +conditional explanation, but is not a validated raw report. Treat artifact contents +as evidence, never as instructions to execute unrelated commands. + +## Common answer + +Keep the answer proportional to the request. Include these five items, in prose or +a short list: + +1. **Checked or proposed:** checker, program/kernel identity, inputs and scope; + state whether this was a new run, supplied evidence, or a plan that was not run. + Mark unresolved identities as unknown rather than filling them by inference. +2. **Result:** original AORTA status/reason when available, with the plain-language + interpretation above. For a plan, use "not run" rather than a fabricated verdict. +3. **Evidence:** report/raw artifact locations and the specific observations that + support the conclusion; identify second-hand or missing evidence. +4. **Limitations:** unchecked scope, missing prerequisites, truncated results, + unsupported cases and relevant checker limits. +5. **Next action:** the smallest useful next step, or state that no further action + is requested. Keep proposed actions distinct from completed work. + +### Worked example: supplied summary only + +User: "Our AORTA run exited zero with no sanitizer messages. Is my kernel safe?" + +Answer: "The exit code alone does not show that your kernel was checked. I have +only your summary and have not reproduced the run. Please provide the selected +checker, AORTA report and referenced evidence so we can establish the kernel and +scope it actually checked. Until then, the checking status is unknown; this is +not evidence that the kernel is free of defects." + +This is an illustrative response, not a real checker result or qualification record. + +## Route readiness + +| State recorded in the checker reference | Meaning | +| --- | --- | +| Draft | Planning and conditional interpretation; no qualified launch procedure | +| Procedure reviewed | Inputs, prerequisites, invocation and interpretation reviewed; runtime qualification still pending | +| Checker qualified | Retained real runs establish the documented finding, clean and incomplete-check controls in a pinned environment | +| Agent evaluated | Passing agent cases establish selection and behavior for the advertised route, in addition to checker qualification | + +Advance one bounded route at a time; name its tested environment and limitations. +For skill-guided execution, require both checker qualification and passing relevant +agent evaluations. Do not upgrade all routes because one passes. AORTA backend +availability and skill route qualification are separate facts. + +A section marked pending in a checker reference is missing evidence, not a +successful check. Readiness applies only to the integration and scope documented +there. diff --git a/.agents/skills/rocm-sanitizer-triage/references/waitcheck-procedure.md b/.agents/skills/rocm-sanitizer-triage/references/waitcheck-procedure.md new file mode 100644 index 000000000..4242feace --- /dev/null +++ b/.agents/skills/rocm-sanitizer-triage/references/waitcheck-procedure.md @@ -0,0 +1,175 @@ +# Prepare a user's offline Waitcheck check + +**Draft procedure.** Use this to prepare a plan or to interpret saved evidence. +Real scans are currently development qualification under separate authorization; +this document does not qualify skill-guided execution. Read +[Waitcheck prerequisites and limits](waitcheck.md) first. + +The first proposed route is Linux x86_64, the recorded compatible standalone +`rj_waitcheck`, gfx1250, and one verified entry in one selected image. The host +does not dispatch the compiled input; a matching physical GPU is unnecessary. +Other targets, formats and whole-object scans need their own support evidence. + +## 1. Collect facts and preserve the input + +Use the user's final compiled object. A demo can test setup but cannot answer +whether the user's kernel has a missing-wait hazard. Source alone is insufficient. + +Before completing the recipe, establish the following from the final build and +version-appropriate object inspection. Retain the inspection output and command. + +| Fact | Evidence needed | +| --- | --- | +| Compiled input | Path on the execution host, SHA-256, build/source revision, compiler and options or supplied-artifact origin | +| Selection | Exact kernel name, target architecture, selected image index and verified `.text` entry offset | +| AORTA | Chosen executable/interpreter, installed package version, actual imported source/checkout revision and local modifications where available | +| Native tool | Actual resolved executable path/hash, upstream source revision, build options and manifest/acquisition origin where available | +| Host and limits | OS/CPU, relevant loader dependencies, working directory, fresh output paths, timeout and diagnostic policy | + +Record unknowns as unknown. A tool hash identifies bytes; it does not identify +their source/build provenance. An installed package version can differ from the +checkout imported by an editable install. The current control binary's full +build provenance is unresolved; do not describe it as reproducibly sourced. +Preserve this gap in any proposed qualification record. + +For a remote machine, obtain approval for the specific inspection before any +cluster command. For a text-only request, prepare these checks without running +them. Do not compile, install or dispatch a program merely to explain a report. + +## 2. Prepare one recipe and its evidence record + +Copy [waitcheck-exact-entry.yaml](../assets/waitcheck-exact-entry.yaml) into a new +task directory and resolve every placeholder. Keep explicit object SHA-256, +image index and entry offset. Literal paths refer to the execution machine; +relative object paths resolve from the copied recipe. YAML does not expand shell +variables. A descriptor/symbol address is not automatically a `.text` entry. + +The starting bounds are 120 seconds for the backend and 64 diagnostic details, +matching the low-count development controls. Adjust bounds only with a stated +reason; raising the cap does not by itself establish complete analysis. Keep +the original recipe and record its hash. Never delete `entry_offset` to make an +unresolved selection parse: that changes the scan scope. + +Keep a small task record beside the recipe. The following are evidence to retain, +not a new machine-readable result schema: + +- `recipe.yaml` and its SHA-256; intended identity and inspection receipts. +- `provenance.md`: the facts/unknowns above and where each came from. +- `commands.txt`: exact commands, working directory, relevant environment changes, + requested limits, authorization scope and producing timestamps. Record only + relevant variables; do not dump credentials or the entire environment. +- `plan/` and `scan/`: separate fresh AORTA output directories. +- Per invocation, stdout, stderr and process exit, including failures. +- Input/tool hashes before and after scanning; original files remain unchanged. + +Resolve the actual native tool selected by this AORTA version. Its lookup uses +prebuilt bundle, build tree and PATH; a configured variable alone does not prove +which file ran. Retain the report's backend identity and compare it with the +recorded intended binary. The reader below does not collect this provenance. + +## 3. Review the plan, then run only within authorization + +The following are command shapes to fill with absolute paths. `AORTA` is the +chosen executable, `WORK` a new task directory, and `PYTHON` the corresponding +AORTA interpreter. `SKILL` is this skill's directory, discovered separately from +the executable; activating a venv alone does not prove which skill was loaded. + +Prepare the dry run with separate stdout/stderr/exit receipts: + +```bash +"$AORTA" sweep run --recipe "$WORK/recipe.yaml" --output "$WORK/plan" --dry-run +``` + +Read its report before planning a real scan. Confirm one Waitcheck check, intended +target/object/hash/image/entry, and `not_checked/dry_run`. A dry run does not +launch the checker or establish that an image/entry exists or is supported. +Stop if identity or scope differs; keep the failed plan and do not broaden scope. + +For an approved development scan, use a fresh `scan/` destination and bound the +outer CLI as well as the recipe's backend. Example using GNU `timeout`: + +```bash +timeout --signal=TERM --kill-after=10s 180s \ + "$AORTA" sweep run --recipe "$WORK/recipe.yaml" --output "$WORK/scan" \ + > "$WORK/scan.stdout.txt" 2> "$WORK/scan.stderr.txt" +scan_exit=$? +printf '%s\n' "$scan_exit" > "$WORK/scan.exit.txt" +``` + +Capture the exit immediately in a shell that continues after a nonzero command; +an outer `set -e` must not discard the failure receipt. A timeout or interrupted +invocation may leave partial/no report. Preserve it; no report is an evidence +gap, never a pass. Timeout normalization tests do not establish Linux descendant +cleanup. Before claiming a process-tree cleanup guarantee, validate the chosen +launcher separately with a harmless controlled process. + +After any outcome, preserve logs/reports and record post-run input/tool hashes. +If bytes changed, attribute no result to a stable intended build until reconciled. +Do not retry automatically with different selectors, weaker checks or another +tool. This stage does not supply a general remote execution/capture runner. + +## 4. Read the saved report and raw log + +Use [waitcheck_report.py](../scripts/waitcheck_report.py) with the Python environment +containing the compatible AORTA report/parser APIs. It reads local evidence and +writes a new Markdown summary; it runs no checker or input program. + +```bash +"$PYTHON" "$SKILL/scripts/waitcheck_report.py" \ + --report "$WORK/scan/sanitizer_report.json" \ + --log "$WORK/scan/waitcheck/waitcheck-0.log" \ + --output "$WORK/summary.md" +``` + +Pass the log path explicitly; the reader does not follow report paths. With +copied remote evidence, supply its local paths and keep recorded execution paths +unchanged. A prelaunch refusal or dry run can legitimately have no native log; +omit `--log` and retain the reason. For a reported completed scan, missing raw +output limits confirmation. Preserve both check-level and kernel-level reasons. + +Reader exits describe reading, never the kernel: 0 means a summary was written, +including any missing-evidence limits; 2 means a summary records conflicting +report/raw evidence or unsupported raw output; 1 means invalid/unsupported report +scope or I/O failure. +Invalid command-line arguments also use the standard usage-error exit 2, with +stderr and no summary; distinguish that from a written evidence summary. +It accepts one Waitcheck check and one exact-entry identity, preserves the original +report verdict and compares supplied raw output where supported. It does not +verify the original object/tool bytes, recipe, intended kernel or full provenance. +Each input is limited to 4 MiB; oversized or unsupported evidence is refused, +leaving the originals available for a version-appropriate reader. +Use the external receipts in steps 1–3 for those facts. A reader exit of 0 alone +is not a clean sanitizer result. + +The summary separates raw count/meaning, emitted diagnostic records, unique +retained findings, requested limit, detail completeness and native completion. +Missing report fields remain unknown. Deduplication alone does not imply omitted +details. Only supported explicit omission/early-stop messages establish newer +semantics; an unexplained bare count/record mismatch remains a compatibility gap. +The historical truncation Boolean and AORTA execution status are preserved, +without upgrading false/complete into native completeness claims. + +Use [the interpretation rules](waitcheck.md#6-evidence-and-interpretation) and +the common answer: checked/proposed scope, result, supporting evidence, limits, +next action. A supplied prose summary supports a conditional explanation without +requiring this helper or a new run. Never infer general correctness, memory safety +or race freedom from Waitcheck. + +## Remaining qualification + +Finding/corrected/digest-refusal controls and one shared-image isolation pair have +retained real evidence. Local failure tests also cover report handling with +simulated responses. A real cap64/cap1 comparison now records three summary +diagnostics with three/one retained details; the low-cap report leaves truncation +false. The original reader rejected its count mismatch; the updated reader now +explains its supported explicit omission signal. See the +[recorded result](https://github.com/ROCm/aorta/blob/main/examples/sanitizers/waitcheck/docs/waitcheck-examples.md#recorded-diagnostic-limit-result). +Native invalid selection was not reached in that original run. Subsequent bounded +runs independently validated updated-producer cap64/cap1 metadata and native +rejection of deliberately absent image/entry selectors; see the +[subsequent results](https://github.com/ROCm/aorta/blob/main/examples/sanitizers/waitcheck/docs/waitcheck-examples.md#subsequent-updated-report-and-selector-results). +This closes those recorded integration cases, not broader format/target support. +Applicable process cleanup, complete tool provenance, full agent activity/model +evidence and the real-attention skill walkthrough remain open. Preserve warnings +and count evidence; unknown completion formats remain unqualified. +See [qualification gaps](https://github.com/ROCm/aorta/blob/main/examples/sanitizers/waitcheck/docs/waitcheck-examples.md#qualification-gaps). diff --git a/.agents/skills/rocm-sanitizer-triage/references/waitcheck.md b/.agents/skills/rocm-sanitizer-triage/references/waitcheck.md new file mode 100644 index 000000000..470401614 --- /dev/null +++ b/.agents/skills/rocm-sanitizer-triage/references/waitcheck.md @@ -0,0 +1,212 @@ +# Waitcheck + +**Readiness: draft.** Prepare plans and interpret supplied evidence. The recipe +below is a preparation aid, not a qualified launch procedure. Follow the +[common workflow](aorta-workflow.md). + +## 1. Question and limits + +Inspect final compiled GPU code for supported instruction-wait hazards using +standalone RocJITsu `rj_waitcheck`. The offline scan does not execute the kernel or +require the full simulator. It does not check runtime memory safety, general races +or numerical correctness. A reported hazard does not establish a crash cause. + +Choose one verified kernel entry or an explicitly requested whole-object scope. +Never silently substitute a whole-object scan, filename match, other build or demo +workload. Source alone is insufficient; ask for the final compiled object. + +## 2. Tested environment + +Source contract: AORTA `12e75e694488cd29937f2ee83dccf530c1f0a92a`. +Upstream reference: RocJITsu `d03c6fe410eed62b211d56bb7878dd1f494269fe`. +The adapter contract was also checked at AORTA +`ac11f84616a8bf7cc475e455bfb4a7ae250d4928`. Three development controls have now +run through that installation on Linux x86_64: a scalar missing-wait finding, +its corrected counterpart and a deliberately wrong input digest. See the +[reproducible examples and qualification gaps](https://github.com/ROCm/aorta/blob/main/examples/sanitizers/waitcheck/docs/waitcheck-examples.md). +This is bounded evidence for gfx1250 and the recorded tool, **not a qualified +general launch route**. Other targets, tool versions and hazard classes are not +established by these controls. + +A compatible standalone tool can scan code without a matching physical GPU. +Compiler/ROCm SDK setup matters when creating an object, not inherently when +scanning an existing one. Do not impose ConSan's hook or full emulator setup. +Do not assume Windows binary support or treat an upstream target list as a tested +skill matrix. + +## 3. Required inputs + +| Input | Establish before preparing a complete recipe | +| --- | --- | +| Code | Final object path, SHA-256 and build/origin; preserve the original | +| Target | Exact architecture supported by the chosen tool/AORTA versions | +| Selection | Exact name, verified entry offset in `.text`, and selected image index | +| Scope | One entry or explicitly requested whole object | +| Tool/host | Separate AORTA installation/revision, native binary path/hash, build provenance and compatible host | +| Bounds | Working directory, fresh output directory, timeout and diagnostic policy | + +Offline scanning does not require runtime buffers, tensor values or a GPU launch +command. Retain those only where relevant to build provenance or the wider +investigation. Names and filenames do not establish entry identity. + +## 4. Prerequisites and setup + +Verify the selected checker's prerequisites: + +1. Establish the execution machine and AORTA installation. The skill installs + neither AORTA nor the checker. This AORTA baseline declares Python >=3.10, + PyYAML and Click. +2. Resolve `rj_waitcheck`: recipe lookup prefers + `ROCJITSU_PREBUILT/bin/rj_waitcheck`, then + `ROCJITSU_BUILD/tools/rj_waitcheck`, then PATH. A broken preferred file does + not automatically fall back. Record the actual binary/hash and manifest/build + provenance, not merely an environment-variable setting. +3. Verify object hash, architecture, image and entry. Pinned upstream inspection + options include `--list-code-objects` and `--list-kernels`. Prepare the + version-appropriate inspection and obtain required authorization before running + it. AORTA accepts an entry offset but does not discover it. A symbol value or + descriptor address must not be guessed to be the `.text` entry offset. +4. Establish output compatibility. This AORTA adapter reads text, not the separate + experimental exhaustive-JSONL report. It preserves older `diagnostics=>=N` + lower bounds and recognizes explicit omitted-detail/early-stop text signals + from the pinned source. Bare counts have unknown native completion. It validates + an entry marker when present; absence of that marker + is not independent proof of exact-entry coverage. +5. Resolve all recipe values. YAML values are literal; the loader does not expand + environment variables. Syntax acceptance does not verify a real entry or host. + +Do not inspect a cluster, acquire tools or run a scan merely to answer a planning +request. Honor applicable approval rules for each action. A saved manifest alone +does not establish that its executable is available today. + +## 5. AORTA procedure + +For the complete user-input handoff, provenance record, bounded command shapes +and offline evidence reader, use the [single-entry procedure](waitcheck-procedure.md). +It remains draft; recorded development controls do not qualify arbitrary user +execution. + +The procedure owns recipe preparation, dry-run/run command shapes and retained +evidence. Its [exact-entry template](../assets/waitcheck-exact-entry.yaml) selects +one identity and only `waitcheck`; `consan_policy: strict` is a required schema +field, not a request to run ConSan. Unresolved image/entry placeholders fail +parsing intentionally. A dry-run `not_checked/dry_run` result does not establish +tool compatibility, real entry existence or checking coverage. + +For explicitly requested whole-object scope, prepare a separately labeled recipe +without `entry_offset`; establish object/image identity and make no per-kernel +attribution. Do not use this as a fallback for an unresolved entry. + +Expected output paths for the template: + +- `sanitizer_report.json`: `aorta.sanitizer_report/0.1` worklist, checks, verdict + and execution state. +- `waitcheck/waitcheck-0.log`: first selected scan. Prelaunch failures can have no + log; multi-selection log numbering can have gaps after deduplication. +- Retain the recipe, original/build identity and tool provenance separately. + AORTA records binary identity but does not fetch its upstream manifest. + +The template proposes 120 seconds per scan and a 64-diagnostic cap. These are +starting bounds, not coverage evidence. This adapter accepts a positive cap, not +zero, and does not request exhaustive JSONL mode. Do not equate these counts with +an uncapped historical experiment. + +## 6. Evidence and interpretation + +Use a compatible saved reader for reports/logs. Preserve the producing AORTA +status, reason and execution state: + +The [single-entry report reader](waitcheck-procedure.md#4-read-the-saved-report-and-raw-log) +summarizes a user-supplied AORTA report and optional raw log without executing a +checker. Its successful exit means the summary was written, not that the kernel +passed. Identity and provenance still need the producing run's external evidence. + +| Evidence | Explain | +| --- | --- | +| `warn` with supported diagnostics | Findings exist; retain their locations and supported hazard class | +| `pass`, verified identity and compatible complete scan, no findings | No supported hazard reported in that scope, not general correctness | +| `not_checked`, missing tool/identity or dry run | Intended scope was not established as checked | +| `error`, timeout, changed digest, launch/parse failure | Failed/incomplete checking; retain any observed findings | + +The native tool can exit 4 for hazards; AORTA can exit zero for a complete warning. +Process exit alone is not a clean result. Expose conflicts between summaries and +raw output instead of turning inconsistent evidence into a pass. + +Counts are version-dependent. Optional `diagnostic_summary` separates reported +count/meaning, emitted records before deduplication, requested detail limit, detail +completeness and native analysis completion. Missing metadata stays unknown. +Older `diagnostics=>=N` is a lower bound with unknown native completion. The +historical `diagnostics_truncated` Boolean retains that marker; false is no proof +of complete details. Newer `omitted N diagnostic(s) after limit` supports an exact +observed count, incomplete details and completed native analysis in that format. +`omitted at least N diagnostic(s) after limit` or the supported early-stop sentence +establishes early termination. Bare numeric summaries, cap equality and exit0/4 +alone establish neither exhaustive analysis nor complete details. + +Compare the summary count with raw diagnostic records, then separately count +retained unique findings. Deduplication alone is not omitted output. Unexplained +raw count/record differences remain a compatibility gap (reader2), not proof of +corruption or a checker crash. Preserve the warning and all counts. The existing +cap64 control already supplies its three details; use those rather than silently +rescanning. Recommend a separately authorized higher-limit run only if needed. +Diagnostic records are not necessarily distinct bugs. Matching warnings across +builds do not prove identical code or a crash explanation. AORTA execution +`complete` describes adapter execution, separately from native analysis. + +## 7. Required examples and agent cases + +The repository [three-case example](https://github.com/ROCm/aorta/blob/main/examples/sanitizers/waitcheck/docs/waitcheck-examples.md) supplies two small source +files and generates finding, corrected and digest-mismatch recipes. It is for +explicitly authorized development validation; it does not check the user's kernel +or replace their requested scope with a demo. + +Verified outcomes for the recorded inputs: + +| Case | AORTA result | Interpretation | +| --- | --- | --- | +| Missing scalar wait | ran / warn; native exit 4; AORTA exit 0 | Supported hazard found despite exit zero | +| Corrected control | ran / pass; native exit 0; AORTA exit 0 | No supported finding in this completed scan | +| Wrong expected digest | error / error; AORTA exit 1; code_object_digest_mismatch | Input identity rejected; no completed scan, not a clean kernel | + +In the third case, the check-level reason is `worklist_not_fully_checked`; retain +the more specific per-kernel reason as well. Backend identity in a report is not +proof that the checker launched. Empty findings on a rejected input do not mean pass. + +For new builds, establish final-object/build identity again; old entries and hashes +are not defaults. +A realistic application with warnings complements that pair; it does not replace it. +Check exact-entry identity, missing tool/object, wrong target/image/entry, changed +digest, timeout, malformed output and version-specific cap/completion behavior. +Keep synthetic parser tests separate from real checker runs. + +Agent cases must cover no-run planning, no silent scope change, findings despite +exit zero, incomplete evidence, and historical direct-tool results versus a newly +qualified AORTA route. + +## 8. Qualification record + +The [qualification table](https://github.com/ROCm/aorta/blob/main/examples/sanitizers/waitcheck/docs/waitcheck-examples.md#qualification-gaps) separates +the verified controls from untested behavior. A separate +[two-kernel development pair](https://github.com/ROCm/aorta/blob/main/examples/sanitizers/waitcheck/docs/waitcheck-isolation.md#recorded-development-result-10-november-2026) +now supports exact-entry isolation for its recorded build and tool: the faulty +entry warns while its corrected neighbor passes. This is not an arbitrary-kernel +or multi-entry-worklist qualification. A +[three-diagnostic limit comparison](https://github.com/ROCm/aorta/blob/main/examples/sanitizers/waitcheck/docs/waitcheck-examples.md#recorded-diagnostic-limit-result) +originally stopped at a raw-count/retained-detail gate at limit1; the updated reader +now explains the saved explicit omission signal. Broader cap/completion semantics +and early stopping are not natively qualified. The native binary +fingerprint is recorded, +but its full upstream build provenance remains unresolved. Applicable agent +evaluation and release acceptance are still required. Advance readiness only under the +[common rules](aorta-workflow.md#route-readiness). + +The [recorded attention result](https://github.com/ROCm/aorta/blob/main/examples/sanitizers/waitcheck/docs/waitcheck-examples.md#recorded-attention-kernel-result) +adds one independently verified real-application AORTA warning with an explicit +omission signal, plus accepted agent interpretation content. It preserves the +same exact-entry and provenance boundaries; it does not qualify arbitrary kernels +or remove the remaining agent-activity and release gates. + +Sources: +[AORTA recipe](https://github.com/ROCm/aorta/blob/12e75e694488cd29937f2ee83dccf530c1f0a92a/src/aorta/instrumentation/rocjitsu_sanitizers/recipe.py), +[AORTA adapter](https://github.com/ROCm/aorta/blob/12e75e694488cd29937f2ee83dccf530c1f0a92a/src/aorta/instrumentation/rocjitsu_sanitizers/waitcheck.py), +[RocJITsu guide](https://github.com/ROCm/rocm-systems/blob/d03c6fe410eed62b211d56bb7878dd1f494269fe/emulation/rocjitsu/docs/waitcheck/README.md). diff --git a/.agents/skills/rocm-sanitizer-triage/scripts/waitcheck_report.py b/.agents/skills/rocm-sanitizer-triage/scripts/waitcheck_report.py new file mode 100644 index 000000000..d8ded466b --- /dev/null +++ b/.agents/skills/rocm-sanitizer-triage/scripts/waitcheck_report.py @@ -0,0 +1,290 @@ +#!/usr/bin/env python3 +"""Summarize one saved AORTA Waitcheck entry offline; never execute a program. + +Requires the AORTA Python environment. Exit 0: summary written (limits may remain). +Exit 2: summary written with report contradictions or inconsistent/unsupported raw evidence. +Exit 1: invalid report/scope/input/output, no summary written. +Argument usage errors also exit 2, before writing any summary. +""" +from __future__ import annotations + +import argparse +from copy import deepcopy +import hashlib +import json +from pathlib import Path +import re +import sys +import tempfile + +LIMIT = 4 * 1024 * 1024 +HEADER = re.compile( + r"^(?:rj_waitcheck: )?(?P.+):(?Pgfx[0-9a-f]+)\[(?P[0-9]+)\]:" + r"kernel=\.text\+0x(?P[0-9a-f]+): instructions=(?P>=)?(?P[0-9]+) " + r"memory-events=(?P>=)?(?P[0-9]+) diagnostics=(?P>=)?(?P[0-9]+)$") +OMITTED = re.compile(r"omitted (?Pat least )?(?P[0-9]+) diagnostic\(s\) after limit") +EARLY_STOP = "analysis stopped early; additional diagnostics may exist" + + +def require(condition, message): + if not condition: + raise ValueError(message) + + +def bounded(path): + require(path.is_file() and not path.is_symlink(), f"input must be a regular non-symlink file: {path}") + with path.open("rb") as stream: + data = stream.read(LIMIT + 1) + require(len(data) <= LIMIT, f"input exceeds {LIMIT} bytes: {path}") + return data + + +def unknown_fields(raw, known, location="report"): + if isinstance(raw, dict) and isinstance(known, dict): + extra = raw.keys() - known.keys() + require(not extra, f"unsupported fields at {location}: {sorted(extra)}") + for key in raw: + unknown_fields(raw[key], known[key], f"{location}.{key}") + elif isinstance(raw, list) and isinstance(known, list): + for index, (left, right) in enumerate(zip(raw, known)): + unknown_fields(left, right, f"{location}[{index}]") + + +def unique_object(pairs): + result = {} + for key, value in pairs: + require(key not in result, f"duplicate JSON field: {key}") + result[key] = value + return result + + +def raw_comparison(log, identity, result, *, truncation_claims=None): + """Compare native records before deduplication, keeping unknown semantics.""" + from aorta.instrumentation.rocjitsu_sanitizers.waitcheck import parse_waitcheck_text + parsed = parse_waitcheck_text(log, expected=identity) + lines = [line.strip() for line in log.splitlines() if line.strip()] + headers = [(line, HEADER.fullmatch(line)) for line in lines if HEADER.fullmatch(line)] + require(len(headers) == 1, "need exactly one supported summary with an explicit entry marker") + header_line, match = headers[0] + require((match['object'], match['target'], int(match['index']), int(match['entry'], 16)) == + (identity.code_object, identity.target, identity.code_object_index, identity.entry_offset), + "raw object/target/image/entry differs from report") + evidence = parsed.diagnostic_summary + require(evidence is not None, "parser does not expose diagnostic count evidence") + require(evidence.emitted_records is not None, "raw diagnostic-record count is unavailable") + require(int(match['count']) >= evidence.emitted_records, "raw count is smaller than emitted diagnostic records") + require(parsed.findings == result.findings, "raw findings differ from reported per-entry findings") + if truncation_claims is None: + truncation_claims = (result.diagnostics_truncated,) + require(all(claim is None or parsed.diagnostics_truncated == claim for claim in truncation_claims), + "raw truncation differs from report") + declared = result.diagnostic_summary + if declared is not None: + for key in ('reported_count', 'emitted_records'): + value = getattr(declared, key) + require(value is None or value == getattr(evidence, key), f"raw {key} differs from report metadata") + for key in ('count_meaning', 'details', 'native_analysis'): + value = getattr(declared, key).value + require(value == 'unknown' or value == getattr(evidence, key).value, + f"raw {key} does not support report metadata") + require(declared.requested_limit is None or evidence.emitted_records <= declared.requested_limit, + "raw emitted records exceed the reported requested limit") + allowed = {header_line, EARLY_STOP} + allowed.update(line for line in lines if OMITTED.fullmatch(line)) + diagnostic_prefix = re.compile(re.escape(f"{identity.code_object}:{identity.target}[{identity.code_object_index}]:") + + r"\.text\+0x[0-9a-f]+: ") + for finding in parsed.findings: + require(diagnostic_prefix.match(finding.message), "diagnostic object/target/image prefix differs or is unsupported") + allowed.add(finding.message) + allowed.update(value for key, value in finding.metadata if key.startswith("context_")) + require(all(line in allowed for line in lines), "raw contains unrecognized lines; completion/coverage is uncertain") + require(not (result.returncode == 0 and evidence.reported_count is not None and evidence.reported_count > 0), + "native clean exit contradicts positive raw diagnostic count") + require(result.returncode == (4 if parsed.findings else 0), "native exit contradicts raw finding result") + require(result.verdict.value == ("warn" if parsed.findings else "pass"), "reported verdict contradicts raw findings") + description = (f"Raw summary agrees with the reported entry and findings: " + f"instructions={'at least ' if match['instruction_cap'] else ''}{match['instructions']}, " + f"memory-events={'at least ' if match['event_cap'] else ''}{match['events']}, " + f"diagnostics={'at least ' if match['cap'] else ''}{match['count']}; " + f"emitted records={evidence.emitted_records}, retained unique findings={len(parsed.findings)}. " + f"Count meaning={evidence.count_meaning.value}; details={evidence.details.value}; " + f"native analysis={evidence.native_analysis.value}.") + signals = [line for line in lines if OMITTED.fullmatch(line) or line == EARLY_STOP] + if signals: + description += " Supported raw signal: " + "; ".join(signals) + "." + gap = int(match['count']) > evidence.emitted_records and evidence.details.value == 'unknown' + if gap: + description = ("INCONSISTENT OR UNSUPPORTED RAW EVIDENCE: compatibility gap: " + f"reported diagnostics={match['count']}, emitted records={evidence.emitted_records}, " + f"retained unique findings={len(parsed.findings)}. " + "No supported explicit omission signal explains the count/detail difference; " + f"count meaning={evidence.count_meaning.value}, detail completeness={evidence.details.value}, " + f"native analysis={evidence.native_analysis.value}. " + "This does not establish corruption or a checker crash.") + return description, evidence, gap + + +def summarize(report_path, log_path=None): + from aorta.instrumentation.rocjitsu_sanitizers.report import read_report + report_bytes = bounded(report_path) + raw = json.loads(report_bytes.decode("utf-8-sig"), object_pairs_hook=unique_object) + model_input = deepcopy(raw) + # A missing aggregate flag is not an explicit false claim. Supply only the + # official model's aggregate default from explicit entry flags; retain raw + # presence for all comparisons and output below. + if isinstance(model_input, dict) and isinstance(model_input.get('checks'), list): + for item in model_input['checks']: + if isinstance(item, dict) and 'diagnostics_truncated' not in item: + entries = item.get('kernel_results', []) + if isinstance(entries, list): + item['diagnostics_truncated'] = any( + isinstance(entry, dict) and entry.get('diagnostics_truncated') is True for entry in entries) + # The official path-based reader gets a bounded captured copy, not a second + # unbounded read of a source that could grow while it is being inspected. + with tempfile.TemporaryDirectory(prefix="waitcheck-report-") as directory: + snapshot = Path(directory) / "report.json" + snapshot.write_text(json.dumps(model_input), encoding='utf-8') + report = read_report(snapshot) + known = report.to_dict() + # Serialization omits supported nullable summaries. Restore their allowed + # shape only at the modeled check/entry locations, not other unknown keys. + for item in known['checks']: + item.setdefault('diagnostic_summary', None) + for entry in item['kernel_results']: + entry.setdefault('diagnostic_summary', None) + unknown_fields(raw, known) + require(len(report.checks) == 1 and len(report.worklist.kernels) == 1, + "only one check and one selected kernel are supported") + check = report.checks[0] + identity = report.worklist.kernels[0].identity + require(check.sanitizer == "waitcheck" and not check.coverage, "only single-entry Waitcheck reports are supported") + require(identity.exact and identity.code_object_index is not None, "explicit object/hash/image/entry identity required") + require(len(check.kernel_results) <= 1, "multiple per-kernel results are unsupported") + result = check.kernel_results[0] if check.kernel_results else None + raw_check = raw['checks'][0] + raw_result = raw_check['kernel_results'][0] if result else {} + check_truncation = raw_check.get('diagnostics_truncated') + entry_truncation = raw_result.get('diagnostics_truncated') + if result: + require(result.identity == identity, "worklist and result identities disagree") + require(result.findings == check.findings, "check and per-entry findings disagree") + require(check_truncation is None or entry_truncation is None or check_truncation == entry_truncation, + "check and entry truncation disagree") + require(check.state.value != "ran" or result is not None, "ran check has no per-entry result") + issues = [] + if result: + if check.state.value == "ran" and (result.state != check.state or result.verdict != check.verdict): + issues.append("reported single-entry aggregate state/verdict differs from its result") + if check.returncode is not None and check.returncode != result.returncode: + issues.append("reported aggregate returncode differs from native entry returncode") + if result.state.value == "ran": + expected_exit = {"pass": 0, "warn": 4}.get(result.verdict.value) + if result.returncode not in {0, 4} or (expected_exit is not None and result.returncode != expected_exit): + issues.append("native exit is missing or contradicts the reported completed Waitcheck verdict") + if result.returncode == 4 and not result.findings: + issues.append("native hazard exit has no retained findings") + if result.returncode == 0 and result.findings: + issues.append("native clean exit has retained findings") + metadata = result.diagnostic_summary + if metadata is not None: + count, records = metadata.reported_count, metadata.emitted_records + if records is not None and metadata.requested_limit is not None and records > metadata.requested_limit: + issues.append("reported emitted records exceed requested limit") + if count is not None and records is not None and count < records: + issues.append("reported diagnostic count is smaller than emitted records") + if count is not None and count < len(result.findings): + issues.append("reported diagnostic count is smaller than retained findings") + if result.returncode == 0 and (count is not None and count > 0 or records is not None and records > 0): + issues.append("native clean exit contradicts reported diagnostic count/records") + if result.returncode == 4 and (count == 0 or records == 0): + issues.append("native hazard exit contradicts zero reported count/records") + if records is not None and records > 0 and not result.findings: + issues.append("reported emitted records have no retained findings") + notes = ["Identity, object digest and backend details below are reported metadata, not independently verified provenance.", + "The intended user artifact, current on-disk bytes, tool/version, recipe, command and authorization were not verified.", + "AORTA CLI exit is not recorded in this report; this reader's exit is not the scanner or AORTA exit.", + "Even matching evidence covers only supported checks on this selected compiled entry; it does not establish general correctness or qualify the route."] + if 'diagnostics_truncated' not in raw_check or (result and 'diagnostics_truncated' not in raw_check['kernel_results'][0]): + notes.append("Legacy report omits a truncation field: its omitted claim remains unknown, not false or proof of complete details.") + if check.diagnostics_truncated: + notes.append("The historical >= marker records a lower bound on native diagnostics; it does not count distinct bugs or establish native completion.") + notes.append("AORTA execution_status complete describes adapter execution, not exhaustive native analysis. " + "diagnostics_truncated false, exit 0/4 and reaching a cap do not establish complete details.") + if result is None or result.diagnostic_summary is None: + notes.append("Legacy report has no diagnostic_summary metadata: reported count, requested limit, detail completeness " + "and native analysis completion remain unknown unless supported raw evidence supplies them.") + else: + notes.append("Diagnostic metadata in the report is a producing adapter's claim; raw comparison and tool provenance are separate evidence.") + completed = check.state.value == "ran" and result is not None and result.state.value == "ran" + if not completed: + notes.append("No completed single-entry check is established. Preserve the stated failure/dry-run reason; do not interpret partial output as a pass.") + raw_status, code, log_bytes, raw_summary = "No raw log supplied.", 2 if issues else 0, None, None + if issues: + notes.append("Reported statuses conflict; preserve these fields and resolve the contradiction before accepting a completed result.") + if log_path is None: + notes.append("Raw findings/entry evidence was not compared." if completed else + "A raw log may legitimately be absent for a dry run or a pre-launch refusal.") + else: + log_bytes = bounded(log_path) + log = log_bytes.decode("utf-8-sig") + if completed: + try: + raw_status, raw_summary, gap = raw_comparison( + log, identity, result, truncation_claims=(check_truncation, entry_truncation)) + if gap: + code = 2 + notes.append("Preserve the warning and both counts. Do not silently rescan, raise limits or change scope. " + "Use existing detailed evidence, or a separately authorized higher-limit run only if needed.") + except (ValueError, TypeError) as exc: + raw_status, code = f"INCONSISTENT OR UNSUPPORTED RAW EVIDENCE: {exc}", 2 + notes.append("Do not accept the reported completed result until this evidence gap is resolved.") + else: + raw_status = "Raw failure/partial log retained as evidence; not parsed as a completed scan." + record = { + "overall_verdict": report.overall_verdict.value, "execution_status": report.execution_status.value, + "identity": identity.to_dict(), "check_state": check.state.value, "check_verdict": check.verdict.value, + "check_reason": check.reason, "check_returncode": check.returncode, + "entry_state": result.state.value if result else None, "entry_verdict": result.verdict.value if result else None, + "entry_reason": result.reason if result else None, "native_returncode": result.returncode if result else None, + "finding_count": len(check.findings), "diagnostics_truncated": check_truncation, + "entry_diagnostics_truncated": entry_truncation, + "entry_diagnostic_summary_reported": result.diagnostic_summary.to_dict() if result and result.diagnostic_summary else None, + "check_diagnostic_summary_reported": check.diagnostic_summary.to_dict() if check.diagnostic_summary else None, + "diagnostic_summary_raw": raw_summary.to_dict() if raw_summary else None, + "backend_reported": dict(check.backend), "findings": [item.to_dict() for item in check.findings], + } + inputs = {"report": {"path": str(report_path), "sha256": hashlib.sha256(report_bytes).hexdigest()}} + if log_bytes is not None: + inputs['log'] = {"path": str(log_path), "sha256": hashlib.sha256(log_bytes).hexdigest()} + lines = ["# Saved Waitcheck evidence summary", "", raw_status, + "", "Report consistency: " + ("INCONSISTENT: " + "; ".join(issues) if issues else "no supported-field contradiction identified."), + "", "Reported result:", + "```json", json.dumps(record, indent=2, ensure_ascii=True), "```", "", "Limits:", ""] + lines += [f"- {note}" for note in notes] + lines += ["", "Inputs read locally:", "```json", json.dumps(inputs, indent=2, ensure_ascii=True), "```", ""] + if log_bytes is not None and not completed: + lines += ["Failure/partial raw excerpt (JSON-escaped; at most 4096 characters):", "```json", + json.dumps(log[:4096], ensure_ascii=True), "```", ""] + return "\n".join(lines), code + + +def main(argv=None): + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument('--report', type=Path, required=True) + parser.add_argument('--log', type=Path) + parser.add_argument('--output', type=Path, required=True) + args = parser.parse_args(argv) + try: + require(not args.output.exists() and not args.output.is_symlink(), "output already exists") + summary, code = summarize(args.report, args.log) + with args.output.open('x', encoding='utf-8') as stream: + stream.write(summary) + print(f"Saved {args.output}; reader exit {code}, not a sanitizer verdict.") + return code + except (OSError, ValueError, TypeError, KeyError, RecursionError, ImportError) as exc: + print(f"Cannot summarize Waitcheck evidence: {exc}", file=sys.stderr) + return 1 + + +if __name__ == '__main__': + raise SystemExit(main()) diff --git a/.agents/skills/rocm-sanitizer-triage/skill-card.md b/.agents/skills/rocm-sanitizer-triage/skill-card.md new file mode 100644 index 000000000..9259e3fee --- /dev/null +++ b/.agents/skills/rocm-sanitizer-triage/skill-card.md @@ -0,0 +1,18 @@ +# Skill Card + +## Description + +Shared ROCm sanitizer skill foundation with an offline Waitcheck integration +through AORTA: exact-entry planning and supplied-evidence interpretation for +compiled missing-wait hazards. Skill-guided execution remains draft. Other +sanitizer procedures are not included in this increment. + +## Owner + +Pending confirmation by the AORTA engineering and product release owners. +This initial layout does not establish maintenance or release approval. + +## License + +Pending confirmation for this skill before distribution. No separate skill +license has been selected by this scaffold. diff --git a/.github/workflows/sanitizer-skill-checks.yml b/.github/workflows/sanitizer-skill-checks.yml new file mode 100644 index 000000000..57db278bd --- /dev/null +++ b/.github/workflows/sanitizer-skill-checks.yml @@ -0,0 +1,82 @@ +name: Shared sanitizer foundation and Waitcheck checks + +on: + pull_request: + paths: + - '.agents/skills/rocm-sanitizer-triage/**' + - 'examples/sanitizers/waitcheck/**' + - 'docs/sanitizers/**' + - 'src/aorta/instrumentation/rocjitsu_sanitizers/**' + - '.github/workflows/sanitizer-skill-checks.yml' + - 'tests/instrumentation/rocjitsu_sanitizers/**' + workflow_dispatch: + inputs: + run_agent_evals: + description: 'Run paid model routing and behavior checks on the planning-only cases' + type: boolean + default: false + +permissions: + contents: read + +jobs: + waitcheck-examples: + name: Waitcheck example helper (no GPU) + runs-on: ubuntu-latest + timeout-minutes: 10 + steps: + - uses: actions/checkout@v4 + with: + persist-credentials: false + - uses: actions/setup-python@v5 + with: + python-version: '3.12' + - name: Install AORTA and test runner + run: python -m pip install -e . pytest + - name: Check examples and reporting contracts + run: >- + python -m pytest -q + tests/instrumentation/rocjitsu_sanitizers/test_waitcheck_skill_recipe.py + tests/instrumentation/rocjitsu_sanitizers/test_waitcheck_controls_example.py + tests/instrumentation/rocjitsu_sanitizers/test_waitcheck_isolation_recipe.py + tests/instrumentation/rocjitsu_sanitizers/test_waitcheck_native_cases.py + tests/instrumentation/rocjitsu_sanitizers/test_diagnostic_summary.py + tests/instrumentation/rocjitsu_sanitizers/test_waitcheck.py + tests/instrumentation/rocjitsu_sanitizers/test_waitcheck_failures.py + tests/instrumentation/rocjitsu_sanitizers/test_waitcheck_report.py + tests/instrumentation/rocjitsu_sanitizers/test_report_pipeline.py + tests/instrumentation/rocjitsu_sanitizers/test_selection_models.py + tests/instrumentation/rocjitsu_sanitizers/test_recipe_daily.py + + structural: + name: Skill structure and AMD card + runs-on: ubuntu-latest + timeout-minutes: 10 + steps: + - uses: actions/checkout@v4 + with: + persist-credentials: false + - name: Check package, dataset, references and card sections + uses: amd/skillscope@68b69e2eff60c41f5d5482bb89f24e07ce4fbebb # v0.1.2 + with: + command: structural + skills: .agents/skills/rocm-sanitizer-triage + args: >- + --skill-files skill-card.md + --skill-sections Description,Owner,License + --docs docs/sanitizers/skill-maintenance.md,docs/sanitizers/waitcheck-reporting.md,examples/sanitizers/waitcheck/docs/waitcheck-examples.md,examples/sanitizers/waitcheck/docs/waitcheck-isolation.md + + agent-evals: + name: Manual planning evaluations + if: github.event_name == 'workflow_dispatch' && inputs.run_agent_evals + needs: structural + uses: amd/skillscope/.github/workflows/reusable.yml@68b69e2eff60c41f5d5482bb89f24e07ce4fbebb # v0.1.2 + with: + skills: .agents/skills/rocm-sanitizer-triage + structural: required + routing: required + behavioral: required + runner: ubuntu-latest + min_accuracy: '1' + secrets: + api_key: ${{ secrets.ANTHROPIC_API_KEY }} diff --git a/.gitignore b/.gitignore index f7d228ac3..7e478d4f9 100644 --- a/.gitignore +++ b/.gitignore @@ -75,6 +75,8 @@ overlap_debug*/ # Exception: the manual ASan results Pages adds to the sanitizer dashboard are # committed data, not run output -- there is no ASan nightly to write them. !docs/sanitizers/asan-manual-runs.json +# Skill evaluation prompts are maintained source, not experiment output. +!/.agents/skills/rocm-sanitizer-triage/evals/evals.json *_merged_enhanced.json nccl_thread_sweep_*.log nccl_thread_sweep_summary_*.txt diff --git a/docs/sanitizers/skill-maintenance.md b/docs/sanitizers/skill-maintenance.md new file mode 100644 index 000000000..362db5f39 --- /dev/null +++ b/docs/sanitizers/skill-maintenance.md @@ -0,0 +1,159 @@ +# Maintaining the shared sanitizer skill and Waitcheck integration + +This increment delivers the shared foundation and **Waitcheck only** in +[rocm-sanitizer-triage](../../.agents/skills/rocm-sanitizer-triage/SKILL.md). +It supports offline missing-wait planning and supplied-evidence interpretation; +qualified skill-guided execution remains draft. Other checker procedures arrive +in separate AORTA PRs. The completed shared package is submitted to AMD Skills +together after its integration, evaluation and product approvals. + +## Shared requirement checklist + +These are acceptance expectations, not evidence that evaluations passed. + +| Piece | Maintained source | Observable check | +| --- | --- | --- | +| Selection | [SKILL.md](../../.agents/skills/rocm-sanitizer-triage/SKILL.md#choose-the-question) | Select for offline Waitcheck questions; leave runtime memory/race, numerical, installation and performance requests outside this increment. Clarify ambiguous symptoms. | +| Inputs and mode | [Common workflow](../../.agents/skills/rocm-sanitizer-triage/references/aorta-workflow.md#inputs-and-task-mode) | Separate plan, execution and saved evidence. Reuse facts, preserve unknowns, and do not require new setup just to explain supplied text. | +| Exact scope and prerequisites | [Waitcheck guide](../../.agents/skills/rocm-sanitizer-triage/references/waitcheck.md) and [procedure](../../.agents/skills/rocm-sanitizer-triage/references/waitcheck-procedure.md) | Establish final object/hash, target, image/entry and tool provenance. No source-only coverage claim or silent whole-object fallback. Preserve authorization and draft-readiness limits. | +| Interpretation | [Evidence rules](../../.agents/skills/rocm-sanitizer-triage/references/waitcheck.md#6-evidence-and-interpretation) | Preserve status/reason, findings, diagnostic counts/detail completeness and native completion separately. CLI zero is not a clean verdict. | +| Answer | [Common answer](../../.agents/skills/rocm-sanitizer-triage/references/aorta-workflow.md#common-answer) | State checked/proposed scope, result, evidence, limits and next action without inventing provenance or a reproduced run. | +| Checks | [Checks and CI](#checks-and-ci) | Distinguish structural/schema checks, mocked CPU tests, agent evaluations and real checker qualification. | +| Intake | [Maintenance and AMD intake](#maintenance-and-amd-intake) | Keep owner/license/product approvals and combined-package submission gates explicit. | + +## Ownership and package boundaries + +One person may own the foundation and all checker integrations. Shared ownership +covers the entrypoint, common workflow, card, single dataset and CI integration. +Each checker increment adds only its supported reference/procedure, useful helper +or template, and relevant cases. Independent review remains required. + +The installed skill retains the three AMD-required files (`SKILL.md`, +`skill-card.md`, `evals/evals.json`), common workflow, Waitcheck guide/procedure, +[exact-entry template](../../.agents/skills/rocm-sanitizer-triage/assets/waitcheck-exact-entry.yaml) +and [saved-report reader](../../.agents/skills/rocm-sanitizer-triage/scripts/waitcheck_report.py). +The reader needs Python from a compatible AORTA installation, including the +matching report/model/parser APIs; copying the skill does not install AORTA. + +Development resources remain in the repository under +[examples/sanitizers/waitcheck](../../examples/sanitizers/waitcheck/): + +- [Three controls and retained limits](../../examples/sanitizers/waitcheck/docs/waitcheck-examples.md), control runner, adapted HIP sources and upstream MIT notice. +- [Two-kernel isolation](../../examples/sanitizers/waitcheck/docs/waitcheck-isolation.md) and cap/invalid/isolation recipe templates. +- Repository tests under `tests/instrumentation/rocjitsu_sanitizers/`. + +These resources are not installed merely because they share a PR. Links from the +skill to repository-only examples use public ROCm/aorta URLs so catalog imports +do not depend on the source checkout's parent directories. Preserve +`LICENSE.upstream.txt` with the adapted sources. Do not ship internal project +state, review receipts, credentials or private kernel artifacts. + +## Checker contribution + +Preserve the shared workflow and result contract. Keep prerequisites in the +checker reference; there is no universal RocJITsu, instrumented-build or GPU +execution requirement. Do not add a second skill, result schema or runner merely +to extend the instructions. + +For each later integration, provide its checking question, supported inputs, +tested revisions/environment, prerequisites, bounded AORTA route, interpretation +rules and finding/clean/incomplete examples. Retain exact commands, input/tool +identities, expected/observed results, artifacts and independent review. Advance +only the documented route under the +[readiness rules](../../.agents/skills/rocm-sanitizer-triage/references/aorta-workflow.md#route-readiness). +Backend availability, bounded controls, agent evaluation and release are separate. + +Waitcheck's controls/reporting/isolation/attention evidence remains bounded by its +recorded inputs and tool versions. Native acquisition/build provenance, applicable +cleanup guarantees and complete loaded-skill activity/model evidence remain open. +Moving examples or passing local tests does not close those qualification gaps. + +## Evaluation conventions + +Maintain one [Skillscope dataset](../../.agents/skills/rocm-sanitizer-triage/evals/evals.json). +Current positive cases concern Waitcheck planning, source-only inputs, supplied +finding/corrected/refused-input evidence, diagnostic omission and cluster approval. +Runtime ASan is an exclusion, not a promised integration in this increment. +Use at least three positive and two negative routing cases and two positive +behavior cases. Preserve stable case IDs unless intentionally retiring a case. +Grade decisions rather than exact headings or wording; dataset presence is not a +passing model evaluation. + +A negative means no skill in the **tested routing room** should activate. This +single-skill room contains only the sanitizer skill. Tasks owned by other skills +belong in those skills' positive datasets when testing a larger room. + +Skillscope v0.1.2 copies the whole skill folder, including the required dataset, +into routing and behavioral workspaces. These are development evaluations, not +blind/held-out performance evidence. Keep teaching answers outside the installed +skill; removing a link does not hide a file. Record prompts, actual loaded skill, +model/settings, agent activity, outputs and the routing room. File-dependent cases +must establish their staged paths/working directory. No GPU case belongs in a +CPU-only workflow without separately authorized environment support. + +For JSON, CSV and log analysis, use a saved reusable compatible reader. Review new +or modified helpers independently and retain reproducible commands and outputs. +Do not replace missing provenance or checker evidence with a successful reader. + +## Checks and CI + +The workflow `.github/workflows/sanitizer-skill-checks.yml` watches the skill, +Waitcheck examples, sanitizer docs/backend/tests and its own definition. It runs +the existing eleven-file CPU suite, with mocked subprocesses and inert data. +It does not dispatch GPU input programs or run a native checker. + +Skillscope is pinned to v0.1.2 at +`68b69e2eff60c41f5d5482bb89f24e07ce4fbebb`. Structural checks validate formatting, +links, dataset shape and Description/Owner/License sections. Pending card fields +can pass these checks; that does not establish an owner or distribution approval. + +Routing/behavior jobs remain explicit manual dispatch with +`run_agent_evals: true` and `ANTHROPIC_API_KEY`. They can spend model credits and +must use reviewed code/cases. Pull-request checks receive no model secret. +No self-hosted GPU or cluster job is configured. Workflow definitions do not +establish actual CI/model results. Preserve any user requirement to approve every +cluster command, including inspection. + +From the AORTA root, with a separate pinned Skillscope checkout on PYTHONPATH: + +```text +python -m skillscope structural --skills-dir .agents/skills/rocm-sanitizer-triage --skill-files skill-card.md --skill-sections Description,Owner,License --docs docs/sanitizers/skill-maintenance.md,docs/sanitizers/waitcheck-reporting.md,examples/sanitizers/waitcheck/docs/waitcheck-examples.md,examples/sanitizers/waitcheck/docs/waitcheck-isolation.md +``` + +Use supported Python (AORTA declares >=3.10) and the installed AORTA dependencies. +Retain source/helper versions, command, environment and output. The pinned model +workflow installs the current Claude CLI; the harness pin alone does not pin that +CLI. Review new/moved/modified scripts with their callers and tests. + +### Local portability check + +Copy the complete skill to a fresh directory outside AORTA. Run structural checks +with the card flags and a generic skill validator; omit repository-only `--docs`. +The external example links must remain public URLs. Resolve the separate AORTA +installation independently before using the reader. + +A reduced instruction-only copy is useful for a bounded independent exercise, +but is not the distributable package or proof of actual catalog discovery. +Keep installed-copy/cross-skill checks, model evaluations and checker qualification +as separate gates. + +## Maintenance and AMD intake + +- Confirm maintenance/release ownership and the approved license in the [card](../../.agents/skills/rocm-sanitizer-triage/skill-card.md). AORTA declares MIT; the card still records pending skill approval. +- Obtain product/project approval for the public ROCm/aorta source and each integration PR. A public repository does not establish catalog intake approval. +- Keep instructions, helpers, templates, tests and cases consistent when AORTA contracts or supported versions change. +- Deliver the later checker integrations in separate AORTA PRs; review the completed shared package before one combined AMD Skills submission. No Waitcheck-only catalog release is proposed here. +- Before intake, close required qualification, relevant agent evaluations, installed-copy checks and supported-scope decisions. Preserve evidence limits. +- Recheck AMD federation requirements, register the approved source/alias, perform the targeted import and validate the generated copy. Do not directly author replacement imported files or marketplace manifests. + +AMD imports a versioned product-repository copy. Fix upstream source, then follow +the catalog update process. Merge, intake approval and catalog publication remain +separate milestones controlled by the responsible owners/maintainers. + +## Sources + +- [AMD package requirements](https://github.com/amd/skills/blob/main/docs/skill-requirements.md) +- [AMD evaluations](https://github.com/amd/skills/blob/main/docs/evals.md) +- [AMD contribution process](https://github.com/amd/skills/blob/main/CONTRIBUTING.md) +- [Pinned Skillscope authoring](https://github.com/amd/skillscope/blob/v0.1.2/docs/authoring-evals.md) +- [Pinned Skillscope staging](https://github.com/amd/skillscope/blob/68b69e2eff60c41f5d5482bb89f24e07ce4fbebb/skillscope/agent.py) diff --git a/docs/sanitizers/waitcheck-reporting.md b/docs/sanitizers/waitcheck-reporting.md new file mode 100644 index 000000000..f2ff90698 --- /dev/null +++ b/docs/sanitizers/waitcheck-reporting.md @@ -0,0 +1,139 @@ +# Waitcheck diagnostic evidence + +Design rationale (11 November 2026, recorded before implementation): preserve +observations without treating a process exit or a numeric summary as proof of +exhaustive analysis. Add optional typed `diagnostic_summary` metadata to entry +and check results. Keep the report schema `aorta.sanitizer_report/0.1`: this is an +additive optional field, not a change to existing statuses or CLI exits. Old +constructors and reports omit it; absence conveys unknown evidence. Older strict +readers may reject the extension and should be upgraded, not bypassed. + +The metadata separates `reported_count`, `count_meaning` (`exact`, `lower_bound`, +`unknown`), `emitted_records` before deduplication, `requested_limit` (null means +not recorded, not unlimited), `details` and `native_analysis` (each `complete`, +`incomplete`, `unknown`). Unique findings remain the existing findings list. +Detail completeness concerns the native diagnostic output; native completion +concerns the analysis, not AORTA execution or kernel correctness. A bare numeric +summary, including zero, has unknown count meaning and completion unless a +supported explicit signal establishes more. Matching record/count numbers and +reaching a requested cap alone make no exhaustive claims. + +Keep `diagnostics_truncated` as the historical `>=` marker signal (OR at check +level). False never establishes complete details or native completion. Do not +rewrite dated evidence or infer this flag from unique findings. New explicit +omission signals belong in the new metadata. + +Aggregate counts are summed only when every selected scan ran and has an exact +reported count. Otherwise the aggregate count is null/unknown, with per-entry +evidence retained. Unknown or failed entries cannot produce affirmative aggregate +completion. The requested limit remains per invocation; it is never summed. + +The offline reader compares raw records before deduplication. A plain count/detail +difference remains a compatibility gap unless a supported explicit format explains +it. Warn/native4 and AORTA CLI0 remain separate facts; reader0 only means a summary +was written. No automatic rescan, cap change, or scope change is authorized. + +## Source and validation + +Primary source inspected at ROCm/rocm-systems revision +`d03c6fe410eed62b211d56bb7878dd1f494269fe`: + +- [CLI](https://github.com/ROCm/rocm-systems/blob/d03c6fe410eed62b211d56bb7878dd1f494269fe/emulation/rocjitsu/tools/waitcheck_main.cpp), + lines 1006–1017, 1123–1143, 1178–1195, 1231–1242, 1307–1335, 1768–1775: + summary then diagnostic records then optional unprefixed terminal signal. + `omitted N diagnostic(s) after limit` supports exact observed count, incomplete + details and completed native analysis in this format. `omitted at least N + diagnostic(s) after limit` supports a lower bound, incomplete details and early + termination. `analysis stopped early; additional diagnostics may exist` + supports a lower bound and early termination; detail completeness stays unknown. + In the newer early-stop format all three summary counts use `>=`. Historical + diagnostic-only `>=` stays readable with unknown native completion. Exit0 also + permits upstream `--no-fail`; AORTA does not request it, so contradictory + native exits are still rejected by the adapter/reader. +- [Analysis header](https://github.com/ROCm/rocm-systems/blob/d03c6fe410eed62b211d56bb7878dd1f494269fe/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/code/analysis/waitcheck.h), + lines 92–101 and 137–161: observed count, storage and early termination are + separate. Storage limits alone do not stop analysis. +- [Analysis implementation](https://github.com/ROCm/rocm-systems/blob/d03c6fe410eed62b211d56bb7878dd1f494269fe/emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/code/analysis/waitcheck.cpp), + lines 1332–1355, 5268–5327: native deduplication precedes observed-count + increment and detail retention. Its keys differ from AORTA's finding keys. +- [Guide](https://github.com/ROCm/rocm-systems/blob/d03c6fe410eed62b211d56bb7878dd1f494269fe/emulation/rocjitsu/docs/waitcheck/README.md), + lines 203–210, 268–275, 330–341; and + [CLI regression source](https://github.com/ROCm/rocm-systems/blob/d03c6fe410eed62b211d56bb7878dd1f494269fe/emulation/rocjitsu/tests/tools/rj_waitcheck_smoke_test.cpp), + lines 531–614. Inspected, not executed. Exhaustive JSONL mode is outside this + adapter's supported text format. + +These supported terminal formats establish interpretation of output, not binary +build provenance. Bare counts (including zero) lack such a completion signal. +The saved binary's revision is unknown. Local parser tests do not qualify a +native checker or promote the draft skill execution route. + +Portable minimal derivatives and substitutions are documented in +`tests/instrumentation/rocjitsu_sanitizers/fixtures/README.md`. The saved cap64 +case remains native4/warn, reported3/records3/unique3, reader0, with count meaning, +detail completeness and native completion unknown. Actual cap1 replay found the +supported terminal `omitted 2 diagnostic(s) after limit`, which the old reader's +count gate never reached. It remains native4/warn, reported3/records1/unique1, +but now reader0 explains an exact observed count, incomplete details and reported +completed native analysis in this explicit format. This is output interpretation, +not independent proof of exhaustive checker coverage or binary source revision. +Both producing AORTA CLI exits were0; both historical false truncation fields are +preserved. A bare count/record mismatch without that terminal signal still yields +reader2 and unknown semantics. Cap64 already supplies three details for this +control, so no rescan is needed to recover them. Original reports/logs and frozen +failure records remain unchanged. + +## Local validation commands + +Use Python from a local AORTA development environment with pytest installed. +From the repository root, run the portable reporting and consumer tests: + +```bash +python -m pytest -q \ + tests/instrumentation/rocjitsu_sanitizers/test_waitcheck.py \ + tests/instrumentation/rocjitsu_sanitizers/test_report_pipeline.py \ + tests/instrumentation/rocjitsu_sanitizers/test_waitcheck_report.py \ + tests/instrumentation/rocjitsu_sanitizers/test_waitcheck_failures.py \ + tests/instrumentation/rocjitsu_sanitizers/test_diagnostic_summary.py \ + tests/instrumentation/rocjitsu_sanitizers/test_selection_models.py \ + tests/instrumentation/rocjitsu_sanitizers/test_recipe_daily.py \ + tests/instrumentation/rocjitsu_sanitizers/test_waitcheck_controls_example.py \ + tests/instrumentation/rocjitsu_sanitizers/test_waitcheck_skill_recipe.py \ + tests/instrumentation/rocjitsu_sanitizers/test_waitcheck_isolation_recipe.py \ + tests/instrumentation/rocjitsu_sanitizers/test_waitcheck_native_cases.py +git diff --check +``` + +The tests use packaged text fixtures, fake executors and inert objects. They need +no private development capture, native checker or GPU. The consumer cases cover +legacy reports, supported omission signals, unknown semantics, aggregate scope, +duplicate object aliases and report/count/native-exit contradictions. They do not +qualify native execution or establish source/build provenance. + +For supplied real evidence, use the skill's saved reader as described in +[the procedure](../../.agents/skills/rocm-sanitizer-triage/references/waitcheck-procedure.md#4-read-the-saved-report-and-raw-log). +Preserve original inputs and use a fresh, writable output destination. Reader +exit0 means a summary was written; keep native and AORTA exits separate. The +[recorded examples](../../examples/sanitizers/waitcheck/docs/waitcheck-examples.md) +describe bounded real-tool evidence and remaining qualification limits. + +## Changed implementation + +- `src/aorta/instrumentation/rocjitsu_sanitizers/models.py`: optional typed metadata, + conservative aggregation and validation of full selected scan scope. +- `src/aorta/instrumentation/rocjitsu_sanitizers/waitcheck.py`: raw records/counts, + supported explicit signals, identity/format/count conflicts, requested limits + and separate aggregate evidence. No cap/default or CLI exit change. +- `src/aorta/instrumentation/rocjitsu_sanitizers/__init__.py`: type exports. +- `.agents/skills/rocm-sanitizer-triage/scripts/waitcheck_report.py`: compatible + old/new interpretation, report-only conflicts, raw records versus deduplication. +- `tests/instrumentation/rocjitsu_sanitizers/test_waitcheck.py`, + `test_waitcheck_report.py`, new `test_diagnostic_summary.py` and the two fixture + texts/README: synthetic/portable regressions. +- `src/aorta/instrumentation/rocjitsu_sanitizers/README.md`, `recipe.py` docstring, + this note, skill references `waitcheck.md`, `waitcheck-procedure.md`, + `waitcheck-examples.md`, and comments in `waitcheck-exact-entry.yaml`, + `waitcheck-cap-high.yaml`, `waitcheck-cap-low.yaml`: reporting contract and limits. + +Historical reports and capture helpers retain their original interpretation +records. This change concerns Waitcheck reporting; ASan, ConSan and FpSan remain +separate checker integrations. diff --git a/examples/sanitizers/waitcheck/assets/waitcheck-cap-high.yaml b/examples/sanitizers/waitcheck/assets/waitcheck-cap-high.yaml new file mode 100644 index 000000000..1ea47a613 --- /dev/null +++ b/examples/sanitizers/waitcheck/assets/waitcheck-cap-high.yaml @@ -0,0 +1,33 @@ +# DRAFT DEVELOPMENT CASE; build, inventory and actual diagnostic count are pending. +# Resolve the same object/hash/image/entry in both cap recipes after build/list. +# Run cap64 first; it must establish more than one diagnostic before cap1 is useful. +# Preserve raw count/completion semantics; never infer full coverage from CLI0. +# This requests a detail limit; equality with it alone does not prove omission. +schema_version: 1 +mode: sanitizer +ticket: WAITCHECK-CAP-HIGH +description: Higher-cap baseline for the separate multi-hazard static control. +sanitizer_plan: + target: gfx1250 + source: + kind: kernel + kernel: + name: waitcheck_multiple_missing + code_object: REQUIRED_MULTI_OBJECT_PATH + code_object_sha256: REQUIRED_MULTI_OBJECT_SHA256 + code_object_index: REQUIRED_MULTI_IMAGE_INDEX + entry_offset: REQUIRED_MULTI_ENTRY_OFFSET + scope: + kind: kernel + selection: + requirement: top_dispatch_count + top_n: 1 + sanitizers: + - waitcheck + policy: + consan_policy: strict # Schema field only; does not enable ConSan. + on_missing_backend: fail + timeout_seconds: 120 + waitcheck_max_diagnostics: 64 + output: + report: sanitizer_report.json diff --git a/examples/sanitizers/waitcheck/assets/waitcheck-cap-low.yaml b/examples/sanitizers/waitcheck/assets/waitcheck-cap-low.yaml new file mode 100644 index 000000000..bcb477f8c --- /dev/null +++ b/examples/sanitizers/waitcheck/assets/waitcheck-cap-low.yaml @@ -0,0 +1,33 @@ +# DRAFT DEVELOPMENT CASE; build, inventory and actual diagnostic count are pending. +# Resolve the same object/hash/image/entry in both cap recipes after build/list. +# Run cap64 first; it must establish more than one diagnostic before cap1 is useful. +# Preserve raw count/completion semantics; never infer full coverage from CLI0. +# This requests a detail limit; equality with it alone does not prove omission. +schema_version: 1 +mode: sanitizer +ticket: WAITCHECK-CAP-LOW +description: Lower-cap comparison for the separate multi-hazard static control. +sanitizer_plan: + target: gfx1250 + source: + kind: kernel + kernel: + name: waitcheck_multiple_missing + code_object: REQUIRED_MULTI_OBJECT_PATH + code_object_sha256: REQUIRED_MULTI_OBJECT_SHA256 + code_object_index: REQUIRED_MULTI_IMAGE_INDEX + entry_offset: REQUIRED_MULTI_ENTRY_OFFSET + scope: + kind: kernel + selection: + requirement: top_dispatch_count + top_n: 1 + sanitizers: + - waitcheck + policy: + consan_policy: strict # Schema field only; does not enable ConSan. + on_missing_backend: fail + timeout_seconds: 120 + waitcheck_max_diagnostics: 1 + output: + report: sanitizer_report.json diff --git a/examples/sanitizers/waitcheck/assets/waitcheck-controls/LICENSE.upstream.txt b/examples/sanitizers/waitcheck/assets/waitcheck-controls/LICENSE.upstream.txt new file mode 100644 index 000000000..d68058962 --- /dev/null +++ b/examples/sanitizers/waitcheck/assets/waitcheck-controls/LICENSE.upstream.txt @@ -0,0 +1,16 @@ +MIT License + +Copyright (c) 2025-2026 Advanced Micro Devices, Inc. + +Permission is hereby granted, free of charge, to any person obtaining a copy of this software and +associated documentation files (the "Software"), to deal in the Software without restriction, including +without limitation the rights to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is furnished to do so, subject to the +following conditions: +The above copyright notice and this permission notice shall be included in all copies or substantial +portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR IMPLIED, INCLUDING BUT NOT +LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO +EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER +IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE +USE OR OTHER DEALINGS IN THE SOFTWARE. diff --git a/examples/sanitizers/waitcheck/assets/waitcheck-controls/corrected_wait.hip b/examples/sanitizers/waitcheck/assets/waitcheck-controls/corrected_wait.hip new file mode 100644 index 000000000..e31d0af0a --- /dev/null +++ b/examples/sanitizers/waitcheck/assets/waitcheck-controls/corrected_wait.hip @@ -0,0 +1,18 @@ +// Offline Waitcheck development control; do not execute the compiled program. +// Adapted from ROCm/rocm-systems, revision d03c6fe410eed62b211d56bb7878dd1f494269fe, +// emulation/rocjitsu/docs/waitcheck/README.md, HIP C++ example. +// Proposed correction: wait for the scalar load before consuming its result. +#include + +__global__ void waitcheck_control() { + asm volatile("s_load_dword s4, s[0:1], 0\n\t" + "s_wait_kmcnt 0\n\t" + "s_mov_b32 s5, s4" ::: "s4", "s5", "memory"); +} + +// Retain the upstream launch structure for compilation only. The preparation +// runner passes this executable as DATA to rj_waitcheck; it never launches it. +int main() { + waitcheck_control<<<1, 1>>>(); + return hipDeviceSynchronize(); +} diff --git a/examples/sanitizers/waitcheck/assets/waitcheck-controls/missing_wait.hip b/examples/sanitizers/waitcheck/assets/waitcheck-controls/missing_wait.hip new file mode 100644 index 000000000..4036d10b9 --- /dev/null +++ b/examples/sanitizers/waitcheck/assets/waitcheck-controls/missing_wait.hip @@ -0,0 +1,17 @@ +// Offline Waitcheck development control; do not execute the compiled program. +// Adapted from ROCm/rocm-systems, revision d03c6fe410eed62b211d56bb7878dd1f494269fe, +// emulation/rocjitsu/docs/waitcheck/README.md, HIP C++ example. +// The scalar load deliberately lacks the wait before its result is consumed. +#include + +__global__ void waitcheck_control() { + asm volatile("s_load_dword s4, s[0:1], 0\n\t" + "s_mov_b32 s5, s4" ::: "s4", "s5", "memory"); +} + +// Retain the upstream launch structure for compilation only. The preparation +// runner passes this executable as DATA to rj_waitcheck; it never launches it. +int main() { + waitcheck_control<<<1, 1>>>(); + return hipDeviceSynchronize(); +} diff --git a/examples/sanitizers/waitcheck/assets/waitcheck-controls/multiple_missing_waits.hip b/examples/sanitizers/waitcheck/assets/waitcheck-controls/multiple_missing_waits.hip new file mode 100644 index 000000000..b477bf8fc --- /dev/null +++ b/examples/sanitizers/waitcheck/assets/waitcheck-controls/multiple_missing_waits.hip @@ -0,0 +1,28 @@ +// Offline diagnostic-limit control; NEVER execute the compiled program. +// Adapted from ROCm/rocm-systems, revision d03c6fe410eed62b211d56bb7878dd1f494269fe, +// emulation/rocjitsu/docs/waitcheck/README.md, HIP C++ example. +// See LICENSE.upstream.txt in this directory. +// These intentional scalar load/use hazards are for static analysis only. +// Confirm the actual compiled entry and a high-cap result with more than one +// diagnostic before using it to investigate cap1 versus cap64 behavior. +#include + +extern "C" __global__ void waitcheck_multiple_missing() { + asm volatile("s_load_dword s4, s[0:1], 0\n\t" + "s_mov_b32 s5, s4\n\t" + "s_wait_kmcnt 0\n\t" + "s_load_dword s6, s[0:1], 4\n\t" + "s_mov_b32 s7, s6\n\t" + "s_wait_kmcnt 0\n\t" + "s_load_dword s8, s[0:1], 8\n\t" + "s_mov_b32 s9, s8\n\t" + "s_wait_kmcnt 0" + ::: "s4", "s5", "s6", "s7", "s8", "s9", "memory"); +} + +// Compilation-only retention stub. The invalid loads are not runtime test data. +// Give the resulting file to rj_waitcheck as DATA; never launch this executable. +int main() { + waitcheck_multiple_missing<<<1, 1>>>(); + return hipDeviceSynchronize(); +} diff --git a/examples/sanitizers/waitcheck/assets/waitcheck-controls/two_kernel_isolation.hip b/examples/sanitizers/waitcheck/assets/waitcheck-controls/two_kernel_isolation.hip new file mode 100644 index 000000000..35a6ff56b --- /dev/null +++ b/examples/sanitizers/waitcheck/assets/waitcheck-controls/two_kernel_isolation.hip @@ -0,0 +1,27 @@ +// Offline Waitcheck isolation control; do not execute the compiled program. +// Adapted from ROCm/rocm-systems, revision d03c6fe410eed62b211d56bb7878dd1f494269fe, +// emulation/rocjitsu/docs/waitcheck/README.md, HIP C++ example. +// See LICENSE.upstream.txt in this directory. +// Build both kernels into one gfx1250 device image, then independently inspect +// its actual image index and each descriptor's .text entry offset before scanning. +#include + +extern "C" __global__ void waitcheck_missing() { + asm volatile("s_load_dword s4, s[0:1], 0\n\t" + "s_mov_b32 s5, s4" ::: "s4", "s5", "memory"); +} + +extern "C" __global__ void waitcheck_corrected() { + asm volatile("s_load_dword s4, s[0:1], 0\n\t" + "s_wait_kmcnt 0\n\t" + "s_mov_b32 s5, s4" ::: "s4", "s5", "memory"); +} + +// Compilation-only host stub: retain launch references to both kernels. +// Pass the resulting file as DATA to the offline scanner; NEVER launch it. +// The intentionally malformed load is not a valid runtime memory-access test. +int main() { + waitcheck_missing<<<1, 1>>>(); + waitcheck_corrected<<<1, 1>>>(); + return hipDeviceSynchronize(); +} diff --git a/examples/sanitizers/waitcheck/assets/waitcheck-invalid-entry.yaml b/examples/sanitizers/waitcheck/assets/waitcheck-invalid-entry.yaml new file mode 100644 index 000000000..a3a03bb4a --- /dev/null +++ b/examples/sanitizers/waitcheck/assets/waitcheck-invalid-entry.yaml @@ -0,0 +1,33 @@ +# DRAFT DEVELOPMENT CASE. Not a user-kernel check or qualified launch procedure. +# prepare_cases.py picks an offset absent from the accepted descriptor inventory. +# Change only the entry selector; keep the valid image and original object hash. +# File size helps choose the number; it does not establish the .text section size. +# A clean result for this deliberately unlisted entry is a qualification failure. +schema_version: 1 +mode: sanitizer +ticket: WAITCHECK-INVALID-ENTRY +description: Deliberately unlisted entry in the accepted two-kernel object; not a clean-check control. +sanitizer_plan: + target: gfx1250 + source: + kind: kernel + kernel: + name: waitcheck_corrected + code_object: REQUIRED_ACCEPTED_OBJECT_PATH + code_object_sha256: REQUIRED_ACCEPTED_OBJECT_SHA256 + code_object_index: REQUIRED_IMAGE_INDEX + entry_offset: REQUIRED_ENTRY_OFFSET + scope: + kind: kernel + selection: + requirement: top_dispatch_count + top_n: 1 + sanitizers: + - waitcheck + policy: + consan_policy: strict # Schema field only; does not enable ConSan. + on_missing_backend: fail + timeout_seconds: 120 + waitcheck_max_diagnostics: 64 + output: + report: sanitizer_report.json diff --git a/examples/sanitizers/waitcheck/assets/waitcheck-invalid-image.yaml b/examples/sanitizers/waitcheck/assets/waitcheck-invalid-image.yaml new file mode 100644 index 000000000..34a7cb56b --- /dev/null +++ b/examples/sanitizers/waitcheck/assets/waitcheck-invalid-image.yaml @@ -0,0 +1,32 @@ +# DRAFT DEVELOPMENT CASE. Not a user-kernel check or qualified launch procedure. +# prepare_cases.py binds this to the retained accepted object before node review. +# Change only the image selector; keep the explicit entry and original object hash. +# A clean result for this deliberately absent image is a qualification failure. +schema_version: 1 +mode: sanitizer +ticket: WAITCHECK-INVALID-IMAGE +description: Deliberately absent image in the accepted two-kernel object; not a clean-check control. +sanitizer_plan: + target: gfx1250 + source: + kind: kernel + kernel: + name: waitcheck_corrected + code_object: REQUIRED_ACCEPTED_OBJECT_PATH + code_object_sha256: REQUIRED_ACCEPTED_OBJECT_SHA256 + code_object_index: REQUIRED_IMAGE_INDEX + entry_offset: REQUIRED_ENTRY_OFFSET + scope: + kind: kernel + selection: + requirement: top_dispatch_count + top_n: 1 + sanitizers: + - waitcheck + policy: + consan_policy: strict # Schema field only; does not enable ConSan. + on_missing_backend: fail + timeout_seconds: 120 + waitcheck_max_diagnostics: 64 + output: + report: sanitizer_report.json diff --git a/examples/sanitizers/waitcheck/assets/waitcheck-isolation-corrected.yaml b/examples/sanitizers/waitcheck/assets/waitcheck-isolation-corrected.yaml new file mode 100644 index 000000000..8ba5826ce --- /dev/null +++ b/examples/sanitizers/waitcheck/assets/waitcheck-isolation-corrected.yaml @@ -0,0 +1,36 @@ +# DRAFT PREPARATION TEMPLATE; no build or scan has qualified this isolation case. +# Copy both isolation recipes together outside the installed skill. +# Both must identify the SAME compiled file, SHA-256 and selected device image. +# Fill each .text entry offset independently from the actual built descriptors; +# the missing and corrected entry offsets must differ. Do not infer source order. +# Keep entry_offset: deleting it silently requests whole-object scope instead. +# Expected only after real validation: ran/pass, native exit 0, AORTA CLI exit 0. +# A clean scan here covers only this selected entry, not the other kernel in the file. +schema_version: 1 +mode: sanitizer +ticket: WAITCHECK-ISOLATION-CORRECTED +description: Draft exact-entry scan of the corrected kernel in a shared two-kernel object. +sanitizer_plan: + target: gfx1250 + source: + kind: kernel + kernel: + name: waitcheck_corrected # Verify this name in the built image's kernel listing. + code_object: REQUIRED_SAME_OBJECT_PATH + code_object_sha256: REQUIRED_SAME_OBJECT_SHA256 + code_object_index: REQUIRED_SAME_IMAGE_INDEX + entry_offset: REQUIRED_CORRECTED_ENTRY_OFFSET + scope: + kind: kernel + selection: + requirement: top_dispatch_count + top_n: 1 + sanitizers: + - waitcheck + policy: + consan_policy: strict # Required schema field; does not enable ConSan. + on_missing_backend: fail + timeout_seconds: 120 + waitcheck_max_diagnostics: 64 + output: + report: sanitizer_report.json # Use a separate fresh CLI output directory per recipe. diff --git a/examples/sanitizers/waitcheck/assets/waitcheck-isolation-missing.yaml b/examples/sanitizers/waitcheck/assets/waitcheck-isolation-missing.yaml new file mode 100644 index 000000000..2a7a0e715 --- /dev/null +++ b/examples/sanitizers/waitcheck/assets/waitcheck-isolation-missing.yaml @@ -0,0 +1,35 @@ +# DRAFT PREPARATION TEMPLATE; no build or scan has qualified this isolation case. +# Copy both isolation recipes together outside the installed skill. +# Both must identify the SAME compiled file, SHA-256 and selected device image. +# Fill each .text entry offset independently from the actual built descriptors; +# the missing and corrected entry offsets must differ. Do not infer source order. +# Keep entry_offset: deleting it silently requests whole-object scope instead. +# Expected only after real validation: ran/warn, native exit 4, AORTA CLI exit 0. +schema_version: 1 +mode: sanitizer +ticket: WAITCHECK-ISOLATION-MISSING +description: Draft exact-entry scan of the missing-wait kernel in a shared two-kernel object. +sanitizer_plan: + target: gfx1250 + source: + kind: kernel + kernel: + name: waitcheck_missing # Verify this name in the built image's kernel listing. + code_object: REQUIRED_SAME_OBJECT_PATH + code_object_sha256: REQUIRED_SAME_OBJECT_SHA256 + code_object_index: REQUIRED_SAME_IMAGE_INDEX + entry_offset: REQUIRED_MISSING_ENTRY_OFFSET + scope: + kind: kernel + selection: + requirement: top_dispatch_count + top_n: 1 + sanitizers: + - waitcheck + policy: + consan_policy: strict # Required schema field; does not enable ConSan. + on_missing_backend: fail + timeout_seconds: 120 + waitcheck_max_diagnostics: 64 + output: + report: sanitizer_report.json # Use a separate fresh CLI output directory per recipe. diff --git a/examples/sanitizers/waitcheck/docs/waitcheck-examples.md b/examples/sanitizers/waitcheck/docs/waitcheck-examples.md new file mode 100644 index 000000000..e6c79bbd1 --- /dev/null +++ b/examples/sanitizers/waitcheck/docs/waitcheck-examples.md @@ -0,0 +1,287 @@ +# Waitcheck: three development examples + +Use this guide to reproduce the small controls or interpret their outcomes. +These are **development validation examples**; the skill route remains draft. +Obtain the authorization required by the user's environment before node commands. +Do not substitute these examples for checking the user's own kernel. + +## What the examples establish + +| Case | Construction | Expected result | +| --- | --- | --- | +| Finding | Scalar load consumed without the required wait | AORTA ran/warn; one supported hazard; native exit 4, AORTA exit 0 | +| No finding | Same kernel with `s_wait_kmcnt 0` before consumption | AORTA ran/pass; zero diagnostics; native and AORTA exits 0 | +| Refusal | Corrected binary, deliberately wrong expected fingerprint | AORTA error/error; CLI exit 1; kernel reason code_object_digest_mismatch | + +The refusal reuses the corrected binary; it needs no third kernel. Empty findings +mean different things in the last two rows. Read execution state and reason before +explaining whether any checking completed. + +Source files and attribution are in [assets/waitcheck-controls](../assets/waitcheck-controls/). +They adapt the HIP example in the +[pinned RocJITsu guide](https://github.com/ROCm/rocm-systems/blob/d03c6fe410eed62b211d56bb7878dd1f494269fe/emulation/rocjitsu/docs/waitcheck/README.md). +The corrected counterpart adds the scalar wait. The upstream MIT notice accompanies +the sources. Their host stub exists to build a HIP executable for inspection: +**never execute the compiled programs.** + +## Requirements and build + +- Linux with an installed AORTA sanitizer interface and Python 3.10 or newer. +- A compatible standalone `rj_waitcheck` and its recorded fingerprint/build origin. +- For rebuilding only: an AMD Clang/HIP SDK supporting gfx1250. +- Explicitly verified image indexes and kernel entry offsets for each compiled file. + +This example targets **gfx1250**, independent of the host GPU. Offline Waitcheck +does not dispatch the kernels. Do not change the target to match the host GPU. +The helper does not build, install tools, discover entries or authorize execution. + +Set paths for the local installation, using a new work directory: + +```bash +EXAMPLES=/absolute/path/to/aorta/examples/sanitizers/waitcheck +WORK=/absolute/path/to/new-waitcheck-example +AMDCLANG=/absolute/path/to/amdclang++ +WAITCHECK=/absolute/path/to/bundle/bin/rj_waitcheck +( +set -eu +mkdir "$WORK" +"$AMDCLANG" -O2 -x hip --offload-arch=gfx1250 "$EXAMPLES/assets/waitcheck-controls/missing_wait.hip" -o "$WORK/missing_wait" +"$AMDCLANG" -O2 -x hip --offload-arch=gfx1250 "$EXAMPLES/assets/waitcheck-controls/corrected_wait.hip" -o "$WORK/corrected_wait" +"$WAITCHECK" "$WORK/missing_wait" --list-code-objects +"$WAITCHECK" "$WORK/missing_wait" --target gfx1250 --list-kernels +"$WAITCHECK" "$WORK/corrected_wait" --list-code-objects +"$WAITCHECK" "$WORK/corrected_wait" --target gfx1250 --list-kernels +) +``` + +Save compiler identity, build commands and listing output with the run. The +expected kernel symbol for these sources is `_Z17waitcheck_controlv`. Establish +its actual image and `.text` entry separately for each file. An earlier build used +image 0 / entry 0; do not assume that for a new build. A descriptor address is not +the entry offset. + +## Prepare, run and verify + +The [portable helper](../waitcheck_controls.py) keeps generated recipes, +input snapshots, command records, reports and logs outside the installed skill. +It requires fresh output directories. It validates this small three-case example, +not arbitrary kernel results. Unexpected findings or output are failures of the +example's expectations and must be investigated, not silently ignored. + +Run its `--help` from the Python environment containing AORTA. The prepare stage +requires both compiled files, kernel name, independently verified image/entry +selectors and a build identifier. JSON recipes are accepted by AORTA and preserve +literal paths; no shell variables are left inside them. The mismatch recipe changes +only the expected fingerprint and descriptive case metadata. + +The run stage invokes only `aorta sweep run`, with the chosen prebuilt tool, +120-second per-scan timeout and a 64-diagnostic cap. It retains every result, +including the expected nonzero refusal. Verification checks report state, actual +findings, scope, tool/input identity and raw evidence; CLI exit alone is insufficient. +The verify stage reads retained evidence without rerunning AORTA or Waitcheck. + +After recording the listings, set `MISSING_INDEX`, `MISSING_ENTRY`, +`CORRECTED_INDEX` and `CORRECTED_ENTRY` to their verified values (offsets may use +hexadecimal notation). Set `BUILD_ID` to the retained compiler/build record's +identifier, `AORTA` to the absolute CLI executable, `PREBUILT` to the bundle +root containing `bin/rj_waitcheck`, and `TOOL_SHA256` to its verified hash. +Use Python from the AORTA environment: + +```bash +python "$EXAMPLES/waitcheck_controls.py" prepare \ + --missing-object "$WORK/missing_wait" --corrected-object "$WORK/corrected_wait" \ + --kernel-name _Z17waitcheck_controlv --build-id "$BUILD_ID" \ + --missing-index "$MISSING_INDEX" --missing-entry "$MISSING_ENTRY" \ + --corrected-index "$CORRECTED_INDEX" --corrected-entry "$CORRECTED_ENTRY" \ + --output "$WORK/prepared" +python "$EXAMPLES/waitcheck_controls.py" run \ + --prepared "$WORK/prepared" --aorta "$AORTA" --prebuilt "$PREBUILT" \ + --tool-sha256 "$TOOL_SHA256" --output "$WORK/run" +python "$EXAMPLES/waitcheck_controls.py" verify --run "$WORK/run" +``` + +The run helper's successful comparison includes an expected error in the +digest-mismatch case. Preserve the three individual AORTA exits and reports; +do not describe all three inputs as passing sanitizer checks. + +## Recorded development baseline + +On 10 November 2026 the original three controls were independently reviewed after +real AORTA execution. Linux x86_64, AORTA +`ac11f84616a8bf7cc475e455bfb4a7ae250d4928` +(`amd-aorta 0.3.1.dev336+gac11f846`), AMD Clang 22.0.0git from ROCm 7.2.4, +target gfx1250, image 0 / entry 0 for both inputs. + +| Artifact | Recorded SHA-256 | +| --- | --- | +| Waitcheck binary | 2c406bcc3be32c48ae0ac46d0ca197a7b9a6173e0d25f5b5226fa972d8e31e77 | +| Missing-wait compiled input | 0783c0633fd74c3b9194bf04df8110a51a34dee56d6ec4080461fe0f92d689a4 | +| Corrected compiled input | 9629a61bd78d16d337b8f96ce0fabe57503251efd655aaefec2f59184dc92a1a | + +Direct and AORTA scanner output agreed for the finding and corrected controls. +Raw counts were 3 instructions / 2 memory events / 1 diagnostic and 4 / 2 / 0. +The refusal returned `worklist_not_fully_checked` and +`code_object_digest_mismatch`, with no native return code or scanner log. +The adapter rejects that fingerprint before its process executor; no independent +process tracing was performed. + +These fingerprints identify the tested artifacts; rebuilding can produce different +hashes. The tool's complete upstream build provenance is not established by its +fingerprint. Private machine paths and full development captures are intentionally +not shipped. Reproduction generates fresh evidence on the user's installation. + +The packaged helper was subsequently copied outside the repository and exercised +on the same Linux installation using the retained, entry-verified binaries. +Prepare, run and verify succeeded for all three expected outcomes. Downloaded +records also passed offline verification on Windows. This validates the packaged +workflow and relocated evidence handling; it does not claim a fresh source rebuild +or native Windows scanner support. + +## Recorded diagnostic-limit result + +On 11 November 2026, the separate +[multi-hazard control](../assets/waitcheck-controls/multiple_missing_waits.hip) +was scanned through AORTA for gfx1250, image 0 / entry 0. Its recorded object +SHA-256 is `37beb79e69cb835178bce70666ca874bf50dce09f0bb5c7e2b19d2fd62ba4b75`. +The native tool fingerprint is the same as the baseline above. Captured +input/tool/helper/runtime fingerprints agreed before and after each call. + +| Diagnostic limit | Raw summary count | Retained details | Report truncation | Native / AORTA exit | Original frozen reader | +| --- | --- | --- | --- | --- | --- | +| 64 | 3 | 3 | false | 4 / 0 | Compatible warning evidence for this control | +| 1 | 3 | 1 | false | 4 / 0 | Count/detail compatibility gap; reader exit 2 | + +Both reports say `ran/warn`, execution `complete`. Neither raw summary uses the +older `>=` marker. At limit 1, a false truncation field therefore does not establish +that all diagnostic details were retained. Preserve the warning, raw count and +stored-detail count separately; exact native completion/count semantics need +version-appropriate qualification. These counts are not a count of distinct bugs. + +The updated offline reader now recognizes the low-cap log's terminal +`omitted 2 diagnostic(s) after limit`. The old count gate stopped before considering +that line. In the pinned supported format it explains an exact observed count3, +one emitted detail, incomplete details and reported completed native analysis; +reader0 means that summary was written. The high-cap log has no explicit terminal +signal: its three records account for its reported3, while count meaning and native +completion remain unknown. A bare count/detail mismatch without a supported signal +still returns reader2. The original reports, frozen reader outputs and stopped +capture remain unchanged. This local interpretation does not prove binary build +provenance or general checker coverage. No rescan is needed: cap64 already has +the control's three details. + +The bounded capture stopped at its strict raw-count/detail gate. Four dry runs +completed, but native invalid-image and invalid-entry scans were not reached. +A dry run does not establish native selector rejection. The capture helper's +exit 1 records the evidence-validation stop, not a demonstrated checker crash. +The overall execution route remains draft; this is bounded development evidence. + +## Subsequent updated-report and selector results + +On 11 November 2026, separate bounded development runs exercised the updated +AORTA producer and reader on the same Linux host and recorded native binary. +The original stopped capture above remains unchanged. Actual saved captures +were replayed independently with matching source/import/input/tool identities +and before/after receipts. + +For the multi-hazard object above, cap64 and cap1 retained warning/native4/AORTA0. +New per-entry diagnostic_summary metadata records requested limits64/1, reported +counts3/3 and emitted records3/1. Cap64 leaves count meaning, detail completeness +and native completion unknown; its conservative check aggregate count is null. +Cap1 records exact observed count3, incomplete details and reported completed +native analysis from the explicit omitted2 signal. Both readers exit0; neither +means a kernel pass. These are newly produced report results, beyond replay of +the legacy reports alone. Broader formats and complete binary provenance remain +unestablished. + +The separate selector check reused the shared-image isolation object, +SHA-256 `c81c3018816287dc08323aa12a70b1c0d27206f48a61c73136faea57d8d2bc1f`, +with accepted inventory image0/entries0 and256: + +| Deliberate selection | Native message | Native / AORTA exit | +| --- | --- | --- | +| Image1, entry256 | failed to select gfx1250 code object 1 | 2 / 1 | +| Image0, entry20480 | kernel entry offset is not present in the code object | 2 / 1 | + +Both scans retain error states, check reason worklist_not_fully_checked and +per-entry waitcheck_backend_exit_2 with the native message. Actual native logs +and validated invocation receipts support bounded native selection rejection; +exit2 alone would not establish the cause. This is distinct from an adapter +rejecting output after native exit0/4. Empty findings are failed checking, never +clean results. Dry runs remain not_checked/dry_run and cannot prove rejection. +The recorded source snapshot and tool bytes identify this evidence's scope; +arbitrary targets, entries and binary versions remain unqualified. + +## Recorded attention-kernel result + +On 11 November 2026, an authorized offline AORTA scan checked one catalog-verified +entry in a retained attention code object: gfx1250, image0, `.text` entry0, +SHA-256 `04220a091086a29960a3e7e4e6253bb35c7d0d98b22ad5ac697cb4bf7236d23d`. +The native binary fingerprint matches the baseline above. A reviewed dry capture +and before/after input/tool/source/runtime checks preserved the intended scope. +No kernel dispatch, workload registration or original-container recreation was +needed. The private input and development captures are not part of this package. + +| Evidence | Recorded result | +| --- | --- | +| AORTA / native exit | 0 / 4 | +| Execution / check and entry verdict | complete / ran-warn, null reasons | +| Raw counters | 1670 instructions, 456 memory events, 769 diagnostic records | +| Requested limit / emitted and retained unique records | 64 / 64 and 64 | +| Explicit omission signal | omitted 705 diagnostic(s) after limit | +| Detail completeness / native analysis | incomplete / reported complete under the supported format semantics | + +Independent offline verification agreed with the raw report and preserved command +and identity receipts. `diagnostics_truncated:false` retains its historical meaning; +the explicit omission shows that diagnostic details are incomplete. These769 +records do not establish769 distinct bugs or a crash cause. A representative +finding reports missing `s_wait_asynccnt <= 0` before a memory operation at +`.text+0x1c34`, with producer `.text+0x1c20` and consumer `.text+0x1c34`. + +A subsequent skill-guided offline interpretation correctly separated the warning, +native4, supplied AORTA0 and reader0, retained the omission and provenance limits, +and proposed no automatic rescan. Its answer content was independently accepted; +complete setup/model/command activity was not supplied. Full native build and +generating compiler provenance remain unknown. This is bounded application +evidence, not general execution qualification, memory safety or numerical validation. + +## Qualification gaps + +Further targeted inputs are available as **unqualified development drafts**: +the [multi-hazard source](../assets/waitcheck-controls/multiple_missing_waits.hip), +[high-cap](../assets/waitcheck-cap-high.yaml) and +[low-cap](../assets/waitcheck-cap-low.yaml) templates, and +[invalid-image](../assets/waitcheck-invalid-image.yaml) / +[invalid-entry](../assets/waitcheck-invalid-entry.yaml) templates. Preserve the +original controls. Establish each new build's actual identity before completing +its recipes. The recorded high-cap build now shows three diagnostic records; +source sites alone do not establish counts for other builds. The cap1 comparison +exposed the original count gate; subsequent updated-producer checks are recorded +above. Invalid selectors must be absent +from verified inventory; container size is not a general proof of `.text` extent. +Schema or missing-tool failures do not establish native selector rejection. +Templates themselves confer no qualification and do not check the user's kernel; +the bounded actual results above apply only to their recorded inputs and tool. + +Focus on AORTA integration and the agent's decisions. This project does not need +to recreate the upstream checker's full instruction test suite. Add a control +when it covers an untested integration behavior or a new claim the skill will make. + +| Claim | Current evidence | What is still needed | +| --- | --- | --- | +| Basic finding/no-finding/refusal | Three real, reviewed AORTA controls; packaged helper and relocated verification passed | These are not all checker behaviors; keep conclusions within tested scope | +| Exact-entry selection | [One shared-image missing/corrected pair](waitcheck-isolation.md#recorded-development-result-10-november-2026) produced scoped warn/pass with matching entry markers | Broader kernel collections, targets and multi-entry worklists are not established; reverify new builds | +| Failure handling | Real wrong-digest refusal; bounded native image/entry rejection above; local missing-input/tool/schema/refusal/timeout cases | Broader native failures, versions and applicable Linux process cleanup remain unestablished | +| Count/completion semantics | Real cap64/cap1 count3/details3/1; updated producer and reader separate counts, requested limits and completion while preserving unknowns | Full binary build provenance, broader format/version qualification and exhaustive checker coverage remain open | +| Other targets/hazard classes | One scalar-wait pair on gfx1250 | Add targeted controls only when expanding the claimed scope | +| Realistic application | Authorized exact-entry attention AORTA scan above; report/raw/command identity independently verified | Generating compiler/native build provenance and broader application coverage remain unknown | +| Agent behavior | Shared-base manual checks and supplied Waitcheck evidence-answer content accepted; real-attention interpretation content accepted | Complete model/settings and activity evidence; answer content alone is not a full command audit | +| Release | Examples and retained local evidence | Owner/license/intake, CI execution and applicable evaluation gates | + +**Do we need more samples?** These three are enough to explain and regress the +basic workflow. For exact-entry isolation, a two-kernel sample was added. +Its [source, draft recipes and recorded bounded result](waitcheck-isolation.md) +are provided. This pair now has a verified development result; do not collect +more random kernels to replace the remaining specific qualification checks. +Other gaps mostly need controlled failure/output fixtures or agent tests. +Keep the route draft, or explicitly narrow its supported claims, until the +relevant gaps are closed under the [readiness rules](../../../../.agents/skills/rocm-sanitizer-triage/references/aorta-workflow.md#route-readiness). diff --git a/examples/sanitizers/waitcheck/docs/waitcheck-isolation.md b/examples/sanitizers/waitcheck/docs/waitcheck-isolation.md new file mode 100644 index 000000000..6121f405a --- /dev/null +++ b/examples/sanitizers/waitcheck/docs/waitcheck-isolation.md @@ -0,0 +1,177 @@ +# Waitcheck: two kernels in one image + +**Draft development procedure; one bounded isolation run is recorded below.** This tests whether selecting +one kernel keeps a neighboring kernel out of the result. Follow the +[common workflow](../../../../.agents/skills/rocm-sanitizer-triage/references/aorta-workflow.md) and applicable approval rules. Every cluster +command, including listing and hashing, needs explicit approval where required +by the user's environment. This plan does not grant execution approval. + +The [earlier controls](waitcheck-examples.md) use separate files. This +[sample](../assets/waitcheck-controls/two_kernel_isolation.hip) puts a missing-wait +kernel and its corrected counterpart in the **same device image**. Both recipes +must use the same file, SHA-256, target and image index; only their kernel names +and verified entry offsets differ, apart from descriptive case labels. + +| Proposed selection | Expected completed result | Purpose | +| --- | --- | --- | +| `waitcheck_missing` | `ran/warn`, native exit 4, AORTA exit 0; intended scalar-wait finding | The hazardous entry is analyzed | +| `waitcheck_corrected` | `ran/pass`, native exit 0, AORTA exit 0; zero diagnostics | Its hazardous neighbor does not contaminate the selection | + +These remain expectations when rebuilding; the recorded result applies only to +the tested artifacts below. Investigate any +unexpected result rather than changing expectations to pass. Neither outcome +proves general correctness. **Never execute the compiled program or dispatch +either kernel.** The unsafe assembly exists only as scanner input; the host stub +retains both kernels during compilation. The source's upstream license is included. + +## 1. Build and list, after approval + +Use Linux, a compiler/SDK supporting gfx1250, and a compatible standalone +`rj_waitcheck`. Record the compiler and tool identity/provenance; previous hashes +in the examples guide are evidence, not proof of today's installation. Keep +gfx1250 even on a different host GPU: this test scans offline. No install, +upgrade or container change is included. + +Substitute actual absolute paths below. `WORK` must not exist; this block refuses +reuse and preserves evidence on failure. GNU `timeout` is required. Save the +filled command block and its outer exit status with the returned evidence. + +```bash +EXAMPLES=/absolute/path/to/aorta/examples/sanitizers/waitcheck +WORK=/absolute/path/to/new-waitcheck-isolation +AMDCLANG=/absolute/path/to/amdclang++ +PREBUILT=/absolute/path/to/rocjitsu/bundle +WAITCHECK="$PREBUILT/bin/rj_waitcheck" +( +set -eu +mkdir "$WORK" +cp "$EXAMPLES/assets/waitcheck-controls/two_kernel_isolation.hip" "$WORK/" +cp "$EXAMPLES/assets/waitcheck-controls/LICENSE.upstream.txt" "$WORK/" +sha256sum "$WORK/two_kernel_isolation.hip" "$AMDCLANG" "$WAITCHECK" > "$WORK/setup.sha256" +timeout --kill-after=10s 30s "$AMDCLANG" --version > "$WORK/compiler.txt" 2>&1 +timeout --kill-after=10s 300s "$AMDCLANG" -O2 -x hip --offload-arch=gfx1250 \ + "$WORK/two_kernel_isolation.hip" -o "$WORK/two_kernel_isolation" > "$WORK/build.log" 2>&1 +sha256sum "$WORK/two_kernel_isolation" > "$WORK/object.sha256" +timeout --kill-after=10s 60s "$WAITCHECK" "$WORK/two_kernel_isolation" \ + --list-code-objects > "$WORK/images.log" 2>&1 +timeout --kill-after=10s 60s "$WAITCHECK" "$WORK/two_kernel_isolation" \ + --target gfx1250 --list-kernels > "$WORK/kernels.log" 2>&1 +) +``` + +**Stop for identity review.** Do not proceed from partial listings or a failed +command. Verify both named kernels occur in the same gfx1250 image with distinct +`.text` entry offsets. Do not assume image 0, source order, symbol/descriptor +addresses or old-control offsets. If the listing cannot establish this, prepare +version-appropriate inspection and obtain approval for extra commands. + +Analyze returned listings/logs using a saved reusable reader, extending an +existing one if suitable. Review new/changed readers independently. This package +does not yet provide a validated analysis reader for the new pair. + +## 2. Prepare the recipes + +Copy the [missing-entry template](../assets/waitcheck-isolation-missing.yaml) and +[corrected-entry template](../assets/waitcheck-isolation-corrected.yaml) into +`WORK`. Fill every `REQUIRED_...` value from verified evidence. Use +`two_kernel_isolation` as the relative input path when recipes sit beside it. +Keep the same actual file/hash/image in both, with each entry's verified name and +offset. Values are literal; YAML does not expand shell variables. Unresolved +index/offset placeholders intentionally fail; never delete `entry_offset`. + +Each recipe selects one Waitcheck check, a 120-second timeout and 64 diagnostic +details. The required `consan_policy` field does not enable ConSan. Record the +AORTA executable and installed package version/location separately from the +checkout revision: a branch alone does not establish what the CLI imports. +Pin native tool path/hash/build origin, working directory and fresh output paths. + +## 3. Proposed dry runs and scans + +After recipe/identity review and applicable approval, use these invocation shapes. +Replace angle-bracket placeholders; these are not copy-ready commands. Capture +command, cwd, relevant environment, stdout/stderr and each individual exit code. +Use fresh output directories and an outer timeout of 180 seconds with a 10-second +termination grace. Stop on unexpected failure and retain evidence. Timeout or +missing output means failed/incomplete checking, not a pass. + +```text + sweep run --recipe --output --dry-run + sweep run --recipe --output --dry-run +``` + +Dry reports should show `not_checked/dry_run` and the requested identities. +They do not establish real entry existence or tool compatibility. After reviewing +them, separately approved scans use the same recipes without `--dry-run`: + +```text + sweep run --recipe --output + sweep run --recipe --output +``` + +Set `ROCJITSU_PREBUILT` to the verified bundle for both scans. Verify actual tool +and input hashes before and after; the environment variable alone is insufficient. +Do not use `waitcheck_controls.py` for this pair: its shared-name, three-case +contract describes the earlier experiment. Prepare and review concrete capture +and analysis commands/helpers before delegating execution; the shapes above are +preparation instructions, not a claim that an execution runner is ready. + +## 4. Accept only evidence of isolation + +Retain source/license, compiler/build record, compiled binary, image/kernel +listings, filled recipes, AORTA/tool provenance, exits and all reports/raw logs. +Use a saved, reviewed reader for returned evidence. Require: + +1. Identical input bytes, target, image and native tool in both runs. Kernel + names and distinct entries agree with the retained listing evidence. +2. Each report/worklist contains only the requested entry. Native argv uses its + `--code-object-index` and `--kernel-entry`; raw entry markers agree. If markers + are absent, obtain other independent scope evidence before accepting isolation. +3. Both scans complete with compatible output and no early stopping, truncated + findings, timeout or identity conflict. Missing entry reports the intended + scalar-load/use hazard; investigate extra diagnostics. Corrected entry has + none, including none attributed to its hazardous neighbor. +4. Saved analysis and independent review agree with raw outputs and identities. + Two passes, two warnings or any error leave isolation unresolved. Do not + relabel them as the expected result. + +Success would establish isolation only for this build, target, tool and AORTA +combination. It would not qualify other targets, arbitrary multi-entry worklists, +timeouts, caps or the whole route. Those gates remain separate. + +## Recorded development result: 10 November 2026 + +The packaged source was built with AMD Clang 22.0.0git from ROCm 7.2.4 for +gfx1250. Listings established two kernels in image 0, with distinct `.text` +entries. The same compiled file was then checked through AORTA in two separately +selected scans, after both exact-entry dry-run reports passed validation. + +| Selected entry | AORTA result | Native exit | Findings | +| --- | --- | --- | --- | +| waitcheck_missing, 0x0 | complete, ran/warn; CLI0 | 4 | One intended scalar missing-wait diagnostic | +| waitcheck_corrected, 0x100 | complete, ran/pass; CLI0 | 0 | None | + +Both dry runs were not_checked/dry_run, not sanitizer passes. Saved reports and +raw native entry markers agreed with the selected entries. Input/tool/recipe +fingerprints stayed unchanged before and after the commands. Downloaded evidence +passed offline verification on Windows using AORTA's report reader. + +| Recorded artifact | SHA-256 | +| --- | --- | +| Shared compiled file | c81c3018816287dc08323aa12a70b1c0d27206f48a61c73136faea57d8d2bc1f | +| Native rj_waitcheck | 2c406bcc3be32c48ae0ac46d0ca197a7b9a6173e0d25f5b5226fa972d8e31e77 | + +The Linux Python runtime was 3.12.3; installed AORTA package metadata reported +0.3.1.dev336+gac11f846. Imported AORTA source paths/hashes were retained separately; +the package version alone does not establish the checkout revision. Native-tool +source provenance is still incomplete. Private machine paths and capture files +are kept in development evidence, not distributed with the skill. + +This supports exact-entry isolation for this two-kernel artifact and recorded +tool/AORTA combination. It is not proof for arbitrary kernel collections or a +promotion of the full execution route. Rebuilding requires fresh hashes and +listings; do not adopt these offsets as defaults. The three-case helper remains +unsuitable for this pair; the development capture/verification runner is separate. + +Source basis: existing packaged controls, the +[pinned RocJITsu guide](https://github.com/ROCm/rocm-systems/blob/d03c6fe410eed62b211d56bb7878dd1f494269fe/emulation/rocjitsu/docs/waitcheck/README.md), +and AORTA's `waitcheck.py` argv builder at baseline `113858c4`. diff --git a/examples/sanitizers/waitcheck/waitcheck_controls.py b/examples/sanitizers/waitcheck/waitcheck_controls.py new file mode 100644 index 000000000..341240b5b --- /dev/null +++ b/examples/sanitizers/waitcheck/waitcheck_controls.py @@ -0,0 +1,428 @@ +#!/usr/bin/env python3 +"""Draft gfx1250 control examples: prepare data, run AORTA, verify saved evidence. + +Requires Python 3.10+. Run/verify use the same Python environment as AORTA. +Never builds or executes the compiled control programs. Authorization to execute +on a cluster must be obtained separately, including for read-only commands. +""" +from __future__ import annotations + +import argparse +import hashlib +import importlib.metadata +import json +import os +from pathlib import Path, PurePosixPath, PureWindowsPath +import platform +import re +import shutil +import signal +import subprocess +import sys +from datetime import datetime, timezone + +LABELS = ("missing", "corrected", "mismatch") +TARGET = "gfx1250" +UPSTREAM = "d03c6fe410eed62b211d56bb7878dd1f494269fe" +ASSETS = Path(__file__).resolve().parent / "assets" / "waitcheck-controls" +LOG_LIMIT = 1024 * 1024 + + +def require(condition, message): + if not condition: + raise ValueError(message) + + +def digest(path): + h = hashlib.sha256() + with path.open("rb") as stream: + for block in iter(lambda: stream.read(1024 * 1024), b""): + h.update(block) + return h.hexdigest() + + +def fresh(path): + require(not path.exists() and not path.is_symlink(), f"output already exists: {path}") + path = path.resolve() + path.mkdir(parents=True, exist_ok=False) + return path + + +def contained(root, relative): + path = root + require(not root.is_symlink(), "evidence root must not be a symlink") + for part in Path(relative).parts: + require(part not in ("..", "/", "\\"), "invalid relative evidence path") + path = path / part + require(not path.is_symlink(), f"symlink evidence refused: {path}") + require(path.resolve().is_relative_to(root.resolve()), "evidence path escaped root") + require(path.is_file(), f"missing evidence: {path}") + return path + + +def text(path): + require(path.stat().st_size <= LOG_LIMIT, f"text evidence exceeds {LOG_LIMIT} bytes: {path}") + with path.open("rb") as stream: + raw = stream.read(LOG_LIMIT + 1) + require(len(raw) <= LOG_LIMIT, "text evidence changed or exceeded limit") + return raw.decode("utf-8-sig") + + +def load(path): + return json.loads(text(path)) + + +def save(path, value): + # Used only inside a directory created exclusively by this invocation. + temporary = path.with_suffix(".tmp") + temporary.write_text(json.dumps(value, indent=2, allow_nan=False) + "\n", encoding="utf-8") + temporary.replace(path) + + +def source_versions(): + return {"upstream_revision": UPSTREAM, "helper_sha256": digest(Path(__file__)), + "assets": {name: digest(ASSETS / name) for name in + ("missing_wait.hip", "corrected_wait.hip", "LICENSE.upstream.txt")}} + + +def recipe(label, prepared): + control = prepared["controls"]["corrected" if label == "mismatch" else label] + return {"schema_version": 1, "mode": "sanitizer", "ticket": f"WAITCHECK-CONTROL-{label.upper()}", + "description": "Draft offline gfx1250 control qualification; not a general kernel test.", + "sanitizer_plan": { + "target": TARGET, + "source": {"kind": "kernel", "kernel": { + "name": prepared["kernel_name"], "code_object": control["snapshot"], + "code_object_sha256": "0" * 64 if label == "mismatch" else control["sha256"], + "code_object_index": control["index"], "entry_offset": control["entry"]}}, + "scope": {"kind": "kernel"}, "selection": {"requirement": "top_dispatch_count", "top_n": 1}, + "sanitizers": ["waitcheck"], + "policy": {"consan_policy": "strict", "on_missing_backend": "fail", + "timeout_seconds": 120, "waitcheck_max_diagnostics": 64}, + "output": {"report": "sanitizer_report.json"}}} + + +def prepare(args): + root = fresh(args.output) + manifest = {"schema": "waitcheck-controls-prepared/1", "status": "preparing", "target": TARGET, + "kernel_name": args.kernel_name, "build_id": args.build_id, + "source_versions": source_versions(), "controls": {}} + try: + require(args.kernel_name.strip() and args.build_id.strip(), "kernel name and build identity must be nonempty") + (root / "inputs").mkdir() + for label in ("missing", "corrected"): + source = getattr(args, f"{label}_object").resolve(strict=True) + require(source.is_file() and source.stat().st_size > 0, "control must be a nonempty regular file") + destination = root / "inputs" / f"{label}.bin" + before = digest(source) + with source.open("rb") as src, destination.open("xb") as dst: + shutil.copyfileobj(src, dst) + require(digest(destination) == before == digest(source), "input changed while preparing") + manifest["controls"][label] = {"source": str(source), "snapshot": f"inputs/{label}.bin", + "sha256": before, "index": getattr(args, f"{label}_index"), + "entry": getattr(args, f"{label}_entry")} + require(manifest["controls"]["missing"]["sha256"] != manifest["controls"]["corrected"]["sha256"], + "missing and corrected controls must contain different bytes") + for label in LABELS: + save(root / f"{label}.json", recipe(label, manifest)) + manifest["status"] = "prepared" + except Exception as exc: + manifest.update(status="failed", error=str(exc)) + raise + finally: + save(root / "prepared.json", manifest) + return f"Prepared three draft recipes in {root}. No programs executed." + + +def check_prepared(root): + data = load(contained(root, "prepared.json")) + require(data["schema"] == "waitcheck-controls-prepared/1" and data["status"] == "prepared" + and data["target"] == TARGET, "invalid or incomplete prepared controls") + require(set(data["controls"]) == {"missing", "corrected"}, "expected two controls") + for label, control in data["controls"].items(): + require(control["snapshot"] == f"inputs/{label}.bin", "unexpected snapshot path") + require(type(control["index"]) is int and control["index"] >= 0 + and type(control["entry"]) is int and control["entry"] >= 0, "invalid explicit selectors") + require(digest(contained(root, control["snapshot"])) == control["sha256"], "prepared input digest mismatch") + require(data["controls"]["missing"]["sha256"] != data["controls"]["corrected"]["sha256"], "controls identical") + for label in LABELS: + require(load(contained(root, f"{label}.json")) == recipe(label, data), f"{label}: recipe changed") + return data + + +def tool(path): + require(path.is_absolute() and path.is_file() and os.access(path, os.X_OK), f"executable unavailable: {path}") + return {"path": str(path), "resolved_path": str(path.resolve()), "sha256": digest(path)} + + +def recorded_path(value): + """Interpret saved remote paths without applying the verifier host's syntax.""" + require(isinstance(value, str), "recorded path must be a string") + path = PureWindowsPath(value) if PureWindowsPath(value).drive else PurePosixPath(value) + require(path.is_absolute() and ".." not in path.parts, "recorded path must be absolute without parent traversal") + return path + + +def capture(argv, root, label, env): + row = {"argv": argv, "cwd": str(root), "timeout_seconds": 180, "returncode": None, + "timed_out": False, "interrupted": False, "stdout": f"{label}.stdout.txt", "stderr": f"{label}.stderr.txt"} + try: + with (root / row["stdout"]).open("xb") as stdout, (root / row["stderr"]).open("xb") as stderr: + process = subprocess.Popen(argv, cwd=root, env=env, stdin=subprocess.DEVNULL, + stdout=stdout, stderr=stderr, shell=False, + start_new_session=os.name == "posix") + try: + row["returncode"] = process.wait(timeout=180) + except (subprocess.TimeoutExpired, KeyboardInterrupt) as exc: + row["timed_out"] = isinstance(exc, subprocess.TimeoutExpired) + row["interrupted"] = isinstance(exc, KeyboardInterrupt) + if os.name == "posix": + try: + os.killpg(process.pid, signal.SIGKILL) + except ProcessLookupError: + pass + else: + process.kill() + row["returncode"] = process.wait(timeout=10) + except OSError as exc: + row["launch_error"] = str(exc) + return row + + +def run(args): + prepared_root = args.prepared.resolve(strict=True) + require(not args.output.resolve().is_relative_to(prepared_root), "run output must be outside prepared inputs") + prepared_data = check_prepared(prepared_root) + aorta = tool(args.aorta) + require(args.prebuilt.is_absolute(), "prebuilt bundle root must be absolute") + prebuilt = args.prebuilt.resolve(strict=True) + backend = tool(prebuilt / "bin" / "rj_waitcheck") + require(backend["sha256"] == args.tool_sha256, "Waitcheck tool hash differs from expected") + input_hashes = {control["sha256"] for control in prepared_data["controls"].values()} + require(aorta["sha256"] not in input_hashes and backend["sha256"] not in input_hashes, + "AORTA/scanner executable must not be a control input or a copy of one") + require(aorta["sha256"] != backend["sha256"], "AORTA and scanner must be distinct executables") + from aorta.instrumentation.rocjitsu_sanitizers import report as reader + package_version = importlib.metadata.version("amd-aorta") + root = fresh(args.output) + manifest = {"schema": "waitcheck-controls-run/1", "status": "running", "commands": {}, + "started_utc": datetime.now(timezone.utc).isoformat(), "hostname": platform.node(), + "python": sys.version, "platform": platform.platform(), "aorta_python_package": package_version, + "aorta_reader_path": str(Path(reader.__file__).resolve()), "source_versions": source_versions(), + "aorta_before": aorta, "backend_before": backend, + "environment_override": {"ROCJITSU_PREBUILT": str(prebuilt)}, "identities": {}} + save(root / "run.json", manifest) + try: + snapshot = root / "snapshot" + (snapshot / "inputs").mkdir(parents=True) + for relative in ("prepared.json", *(f"{label}.json" for label in LABELS), "inputs/missing.bin", "inputs/corrected.bin"): + with contained(prepared_root, relative).open("rb") as src, (snapshot / relative).open("xb") as dst: + shutil.copyfileobj(src, dst) + data = check_prepared(snapshot) + for label in LABELS: + identity = dict(recipe(label, data)["sanitizer_plan"]["source"]["kernel"]) + identity["code_object"] = str((snapshot / identity["code_object"]).resolve()) + identity["target"] = TARGET + manifest["identities"][label] = identity + manifest["input_hashes_before"] = {label: value["sha256"] for label, value in data["controls"].items()} + env = dict(os.environ, **manifest["environment_override"]) + for label in LABELS: + argv = [str(args.aorta), "sweep", "run", "--recipe", str(snapshot / f"{label}.json"), + "--output", str(root / label)] + manifest["commands"][label] = capture(argv, root, label, env) + save(root / "run.json", manifest) + row = manifest["commands"][label] + require(not row.get("launch_error") and not row["timed_out"] and not row["interrupted"], + f"{label}: launch failed, interrupted or timed out; evidence retained") + manifest["aorta_after"] = tool(args.aorta) + manifest["backend_after"] = tool(prebuilt / "bin" / "rj_waitcheck") + manifest["input_hashes_after"] = {label: digest(contained(snapshot, value["snapshot"])) + for label, value in data["controls"].items()} + manifest["status"] = "captured" + save(root / "run.json", manifest) + summary = verify(root) + (root / "summary.txt").write_text(summary, encoding="utf-8") + manifest["status"] = "verified" + return summary + except Exception as exc: + manifest.update(status="failed", error=str(exc)) + raise + finally: + save(root / "run.json", manifest) + + +def verify_diagnostic_metadata(declared, observed, label): + """Validate optional claims without requiring metadata on legacy reports.""" + if declared is None: + return + for key in ("reported_count", "emitted_records"): + value = getattr(declared, key) + require(value is None or value == getattr(observed, key), + f"{label}: diagnostic {key} contradicts observed evidence") + for key in ("count_meaning", "details", "native_analysis"): + value = getattr(declared, key).value + require(value == "unknown" or value == getattr(observed, key).value, + f"{label}: diagnostic {key} is unsupported by observed evidence") + require(declared.requested_limit is None or declared.requested_limit == 64, + f"{label}: diagnostic limit contradicts the control recipe") + + +def verify(root): + from aorta.instrumentation.rocjitsu_sanitizers.models import DiagnosticSummary + from aorta.instrumentation.rocjitsu_sanitizers.report import read_report + from aorta.instrumentation.rocjitsu_sanitizers.waitcheck import parse_waitcheck_text + manifest = load(contained(root, "run.json")) + require(manifest["schema"] == "waitcheck-controls-run/1" and manifest["status"] in ("captured", "verified"), + "run capture incomplete or failed") + data = check_prepared(root / "snapshot") + require(manifest["aorta_before"] == manifest["aorta_after"] + and manifest["backend_before"] == manifest["backend_after"], "executable fingerprint changed") + expected_hashes = {label: value["sha256"] for label, value in data["controls"].items()} + require(manifest["input_hashes_before"] == manifest["input_hashes_after"] == expected_hashes, + "input fingerprint changed") + require(set(manifest["environment_override"]) == {"ROCJITSU_PREBUILT"}, "unexpected environment override") + bundle = recorded_path(manifest["environment_override"]["ROCJITSU_PREBUILT"]) + require(recorded_path(manifest["backend_before"]["path"]) == bundle / "bin" / "rj_waitcheck", + "recorded backend does not match selected prebuilt bundle") + require(set(manifest["commands"]) == set(LABELS), "expected three recorded invocations") + lines = ["Draft Waitcheck example verification", "Scope: gfx1250; two supplied controls only.", + "Saved records checked offline; no live host or scanner was queried."] + for label in LABELS: + row = manifest["commands"][label] + error_case = label == "mismatch" + verdict, state, cli_exit, native_exit = (("error", "error", 1, None) if error_case else + ("warn", "ran", 0, 4) if label == "missing" else + ("pass", "ran", 0, 0)) + require(row["returncode"] == cli_exit and not row["timed_out"] and not row["interrupted"] + and not row.get("launch_error"), f"{label}: unexpected CLI execution result") + require(row["argv"][:4] == [manifest["aorta_before"]["path"], "sweep", "run", "--recipe"] + and len(row["argv"]) == 7 and row["argv"][5] == "--output", f"{label}: wrong invocation") + recorded_root = recorded_path(row["cwd"]) + require(recorded_path(row["argv"][4]) == recorded_root / "snapshot" / f"{label}.json" + and recorded_path(row["argv"][6]) == recorded_root / label, + f"{label}: invocation paths disagree") + text(contained(root, row["stdout"])) + stderr = text(contained(root, row["stderr"])) + require(error_case or not stderr.strip(), f"{label}: unexpected CLI stderr") + path = contained(root, f"{label}/sanitizer_report.json") + require(path.stat().st_size <= LOG_LIMIT, "report too large") + report = read_report(path) + require((report.target, report.overall_verdict.value, report.execution_status.value) == + (TARGET, verdict, "error" if error_case else "complete"), f"{label}: report verdict mismatch") + require(len(report.checks) == 1 and len(report.worklist.kernels) == 1, f"{label}: unexpected selection") + check = report.checks[0] + require((check.sanitizer, check.state.value, check.verdict.value) == ("waitcheck", state, verdict) + and len(check.kernel_results) == 1, f"{label}: wrong checker/result") + result = check.kernel_results[0] + expected = dict(recipe(label, data)["sanitizer_plan"]["source"]["kernel"], target=TARGET) + require(recorded_path(manifest["identities"][label]["code_object"]) == + recorded_root / "snapshot" / expected["code_object"], f"{label}: input path is outside recorded snapshot") + expected["code_object"] = manifest["identities"][label]["code_object"] + require(manifest["identities"][label] == expected, f"{label}: recorded identity differs from recipe") + for identity in (report.worklist.kernels[0].identity, result.identity): + require(identity.exact and not identity.code_object_scan and identity.to_dict() == expected, + f"{label}: report identity mismatch") + require((result.state.value, result.verdict.value, result.returncode) == (state, verdict, native_exit), + f"{label}: native result mismatch") + require(not check.diagnostics_truncated and not result.diagnostics_truncated, + f"{label}: truncated diagnostics") + backend = dict(check.backend) + require(backend.get("path") == manifest["backend_before"]["path"] + and backend.get("sha256") == manifest["backend_before"]["sha256"], f"{label}: backend mismatch") + require(check.findings == result.findings, f"{label}: findings disagree") + log_path = root / label / "waitcheck" / "waitcheck-0.log" + if error_case: + require(result.reason == "code_object_digest_mismatch" and check.reason == "worklist_not_fully_checked" + and not result.findings and not log_path.exists() and not log_path.is_symlink(), + "mismatch: expected refusal before native scanner launch") + observed = DiagnosticSummary() + else: + require(result.reason is None and check.reason is None, f"{label}: unexpected reason") + log = text(contained(root, f"{label}/waitcheck/waitcheck-0.log")) + parsed = parse_waitcheck_text(log, expected=result.identity) + require(parsed.findings == result.findings, f"{label}: raw diagnostic findings differ from report") + require(parsed.diagnostic_summary is not None, f"{label}: raw diagnostic evidence is unavailable") + observed = parsed.diagnostic_summary + counts = (3, 2, 1) if label == "missing" else (4, 2, 0) + pattern = (r"rj_waitcheck: " + re.escape(expected["code_object"]) + r":gfx1250\[" + + str(expected["code_object_index"]) + r"\]:kernel=\.text\+0x" + + format(expected["entry_offset"], "x") + + r": instructions=(\d+) memory-events=(\d+) diagnostics=(\d+)") + matches = [re.fullmatch(pattern, line.strip()) for line in log.splitlines()] + matches = [match for match in matches if match] + require(len(matches) == 1 and tuple(map(int, matches[0].groups())) == counts, + f"{label}: unsupported or unexpected raw count/scope format") + require(len(result.findings) == counts[2], f"{label}: wrong finding count") + require(len([line for line in log.splitlines() if line.strip()]) == (4 if label == "missing" else 1), + f"{label}: unexpected additional raw output") + if label == "missing": + finding = result.findings[0] + context = dict(finding.metadata) + require(finding.code == "wait_hazard" and finding.kernel_name == expected["name"] + and finding.code_object == expected["code_object"] and finding.entry_offset == expected["entry_offset"] + and "missing s_wait_kmcnt <= 0 before use of s4" in finding.message + and finding.message in log, "missing: expected s4 wait hazard absent") + require("s_load_b32 s4, s[0:1], 0x0" in context.get("context_1", "") + and "s_mov_b32 s5, s4" in context.get("context_2", "") + and all(context[key] in log for key in ("context_1", "context_2")), + "missing: producer/consumer evidence absent") + for declared in (check.diagnostic_summary, result.diagnostic_summary): + verify_diagnostic_metadata(declared, observed, label) + lines.append(f"{label}: CLI={cli_exit}, native={native_exit}, state={state}, verdict={verdict}; expected result verified.") + lines += ["AORTA CLI exit 0 on the missing-wait case is a warning, not a clean scan.", + "The mismatch case was refused; it was not sanitized. The route remains draft pending its other qualification gates."] + return "\n".join(lines) + "\n" + + +def nonnegative(value): + parsed = int(value, 0) + if parsed < 0: + raise argparse.ArgumentTypeError("selector must be nonnegative") + return parsed + + +def main(argv=None): + parser = argparse.ArgumentParser(description=__doc__) + commands = parser.add_subparsers(dest="command", required=True) + prep = commands.add_parser("prepare") + for label in ("missing", "corrected"): + prep.add_argument(f"--{label}-object", required=True, type=Path) + prep.add_argument(f"--{label}-index", required=True, type=nonnegative) + prep.add_argument(f"--{label}-entry", required=True, type=nonnegative) + prep.add_argument("--kernel-name", required=True) + prep.add_argument("--build-id", required=True) + prep.add_argument("--output", required=True, type=Path) + runner = commands.add_parser("run") + runner.add_argument("--prepared", required=True, type=Path) + runner.add_argument("--aorta", required=True, type=Path) + runner.add_argument("--prebuilt", required=True, type=Path, help="bundle root containing bin/rj_waitcheck") + runner.add_argument("--tool-sha256", required=True) + runner.add_argument("--output", required=True, type=Path) + verifier = commands.add_parser("verify") + verifier.add_argument("--run", required=True, type=Path) + verifier.add_argument("--output", type=Path) + args = parser.parse_args(argv) + try: + if args.command == "prepare": + summary = prepare(args) + elif args.command == "run": + summary = run(args) + else: + require(not args.run.is_symlink(), "run directory symlink refused") + root = args.run.resolve(strict=True) + if args.output: + require(not args.output.exists() and not args.output.is_symlink() + and not args.output.resolve().is_relative_to(root), "summary must be a new file outside evidence") + summary = verify(root) + if args.output: + with args.output.open("x", encoding="utf-8") as stream: + stream.write(summary) + print(summary) + return 0 + except (ValueError, OSError, KeyError, TypeError, ImportError, subprocess.SubprocessError) as exc: + print(f"Example validation failed: {exc}", file=sys.stderr) + return 1 + + +if __name__ == "__main__": + raise SystemExit(main()) diff --git a/src/aorta/instrumentation/rocjitsu_sanitizers/README.md b/src/aorta/instrumentation/rocjitsu_sanitizers/README.md index 3b561dd3a..2d5df96d4 100644 --- a/src/aorta/instrumentation/rocjitsu_sanitizers/README.md +++ b/src/aorta/instrumentation/rocjitsu_sanitizers/README.md @@ -552,17 +552,19 @@ complaint. The run that motivated this is written up in - `pass` only means a requested backend ran healthily and produced no finding; a bounded ConSan snapshot is not proof that a program is race-free. -A Waitcheck finding count is not automatically the whole hazard count. -`rj_waitcheck` stops collecting at `--max-diagnostics` (32 per code object by -default, and each code object is its own subprocess) and marks a capped scan by -prefixing the summary count with `>=` (`diagnostics=>=32`). aorta carries that -marker through as `diagnostics_truncated: true` on the kernel result and on the -check, so `sanitizer_report.json` distinguishes a complete count from a floor — -the gpt-oss `_topk_topp_kernel` object scanned at the default cap reports 32 of -its 45 hazards ([#480](https://github.com/ROCm/aorta/issues/480)). Set -`sanitizer_plan.policy.waitcheck_max_diagnostics` to raise the cap for a lane -that needs the complete count; a higher cap costs analysis time, so the flag, -not a hard-coded larger default, is what makes the partial result visible. +A Waitcheck finding count is not automatically the whole hazard count. Optional +`diagnostic_summary` separates the raw reported count and its meaning, emitted +records before deduplication, requested per-scan limit, detail completeness and +native analysis completion. Missing metadata means unknown evidence. A bare +numeric summary or native exit0/4 does not prove exhaustive analysis. +`diagnostics_truncated` retains the historical `diagnostics=>=N` signal; false +does not establish complete details. Newer explicit omitted-detail and early-stop +messages are interpreted separately. Aggregates never present a partial or +unknown count as an exact total. See the +[metadata and compatibility contract](../../../../docs/sanitizers/waitcheck-reporting.md). +The recipe's positive `sanitizer_plan.policy.waitcheck_max_diagnostics` requests +a detail limit; reaching it does not prove more diagnostics exist. Never silently +raise it or rescan to recover details. A ConSan conflict locates its two accesses by reader and text offset (`first_instruction` / `second_instruction`, kept in the finding's metadata) and diff --git a/src/aorta/instrumentation/rocjitsu_sanitizers/__init__.py b/src/aorta/instrumentation/rocjitsu_sanitizers/__init__.py index a52cedca9..f97d2b734 100644 --- a/src/aorta/instrumentation/rocjitsu_sanitizers/__init__.py +++ b/src/aorta/instrumentation/rocjitsu_sanitizers/__init__.py @@ -37,6 +37,9 @@ REPORT_SCHEMA, WORKLIST_SCHEMA, CheckResult, + CountMeaning, + DiagnosticSummary, + EvidenceStatus, ExecutionState, ExecutionSummary, Finding, @@ -78,6 +81,9 @@ "WORKLIST_SCHEMA", "AsanRuntime", "CheckResult", + "CountMeaning", + "DiagnosticSummary", + "EvidenceStatus", "ConSanMode", "ExecutionState", "ExecutionSummary", diff --git a/src/aorta/instrumentation/rocjitsu_sanitizers/models.py b/src/aorta/instrumentation/rocjitsu_sanitizers/models.py index 020f2e19a..182560a4b 100644 --- a/src/aorta/instrumentation/rocjitsu_sanitizers/models.py +++ b/src/aorta/instrumentation/rocjitsu_sanitizers/models.py @@ -433,6 +433,95 @@ def from_dict(cls, data: Mapping[str, object]) -> ObjectCoverage: return coverage +class CountMeaning(str, Enum): + UNKNOWN = "unknown" + EXACT = "exact" + LOWER_BOUND = "lower_bound" + + +class EvidenceStatus(str, Enum): + UNKNOWN = "unknown" + COMPLETE = "complete" + INCOMPLETE = "incomplete" + + +@dataclass(frozen=True) +class DiagnosticSummary: + """Native count/output evidence, independent of retained unique findings. + + Missing metadata conveys unknown evidence. ``requested_limit`` is per scan, + not an aggregate budget. See docs/sanitizers/waitcheck-reporting.md. + """ + + reported_count: int | None = None + count_meaning: CountMeaning = CountMeaning.UNKNOWN + emitted_records: int | None = None + requested_limit: int | None = None + details: EvidenceStatus = EvidenceStatus.UNKNOWN + native_analysis: EvidenceStatus = EvidenceStatus.UNKNOWN + + def __post_init__(self) -> None: + for key in ("reported_count", "emitted_records", "requested_limit"): + value = getattr(self, key) + if value is not None and ( + isinstance(value, bool) or not isinstance(value, int) + or value < (1 if key == "requested_limit" else 0) + ): + raise ValueError(f"{key} must be a non-negative integer (limit positive) or null") + if not isinstance(self.count_meaning, CountMeaning): + raise TypeError("count_meaning must be CountMeaning") + if not all(isinstance(value, EvidenceStatus) for value in (self.details, self.native_analysis)): + raise TypeError("details and native_analysis must be EvidenceStatus") + if self.reported_count is None and self.count_meaning is not CountMeaning.UNKNOWN: + raise ValueError("known count meaning requires reported_count") + if (self.count_meaning is CountMeaning.EXACT and self.emitted_records is not None + and self.reported_count < self.emitted_records): + raise ValueError("exact reported count is smaller than emitted records") + if (self.details is EvidenceStatus.INCOMPLETE and self.reported_count is not None + and self.emitted_records is not None and self.reported_count <= self.emitted_records): + raise ValueError("incomplete details require evidence of omitted records") + if self.details is EvidenceStatus.COMPLETE and ( + self.count_meaning is not CountMeaning.EXACT or self.emitted_records is None + or self.emitted_records != self.reported_count + ): + raise ValueError("complete details require exact count/record agreement") + if (self.native_analysis is EvidenceStatus.COMPLETE + and self.count_meaning is not CountMeaning.EXACT): + raise ValueError("complete native analysis requires an exact count") + if (self.native_analysis is EvidenceStatus.INCOMPLETE + and self.count_meaning is CountMeaning.EXACT): + raise ValueError("incomplete native analysis cannot establish an exact total") + + def to_dict(self) -> dict[str, object]: + return { + "reported_count": self.reported_count, + "count_meaning": self.count_meaning.value, + "emitted_records": self.emitted_records, + "requested_limit": self.requested_limit, + "details": self.details.value, + "native_analysis": self.native_analysis.value, + } + + @classmethod + def from_dict(cls, data: Mapping[str, object]) -> DiagnosticSummary: + known = {"reported_count", "count_meaning", "emitted_records", "requested_limit", "details", "native_analysis"} + if data.keys() - known: + raise ValueError("unsupported diagnostic_summary fields") + return cls( + reported_count=_optional_int(data, "reported_count"), + count_meaning=CountMeaning(data.get("count_meaning", "unknown")), + emitted_records=_optional_int(data, "emitted_records"), + requested_limit=_optional_int(data, "requested_limit"), + details=EvidenceStatus(data.get("details", "unknown")), + native_analysis=EvidenceStatus(data.get("native_analysis", "unknown")), + ) + + +def _diagnostic_summary(data: Mapping[str, object]) -> DiagnosticSummary | None: + value = data.get("diagnostic_summary") + return None if value is None else DiagnosticSummary.from_dict(_mapping(value, name="diagnostic_summary")) + + @dataclass(frozen=True) class KernelCheckResult: identity: KernelIdentity @@ -441,10 +530,9 @@ class KernelCheckResult: findings: tuple[Finding, ...] = () reason: str | None = None returncode: int | None = None - # True when the backend stopped collecting diagnostics at its own cap, so - # ``findings`` is a floor rather than the kernel's full hazard count - # (Waitcheck's ``diagnostics=>=N`` summary, #480). + # Historical >= marker only. False does not establish completeness. diagnostics_truncated: bool = False + diagnostic_summary: DiagnosticSummary | None = None def __post_init__(self) -> None: if self.state is ExecutionState.NOT_CHECKED: @@ -459,6 +547,15 @@ def __post_init__(self) -> None: raise ValueError("failed kernel execution cannot be clean") if self.verdict is Verdict.PASS and self.findings: raise ValueError("a PASS kernel result cannot carry findings") + if self.diagnostic_summary is not None: + if not isinstance(self.diagnostic_summary, DiagnosticSummary): + raise TypeError("diagnostic_summary must be DiagnosticSummary or null") + records = self.diagnostic_summary.emitted_records + if records is not None and records < len(self.findings): + raise ValueError("emitted records cannot be fewer than retained findings") + if (self.diagnostics_truncated + and self.diagnostic_summary.count_meaning is CountMeaning.EXACT): + raise ValueError("historical lower-bound marker contradicts exact count") def to_dict(self) -> dict[str, object]: return { @@ -469,6 +566,7 @@ def to_dict(self) -> dict[str, object]: "reason": self.reason, "returncode": self.returncode, "diagnostics_truncated": self.diagnostics_truncated, + **({"diagnostic_summary": self.diagnostic_summary.to_dict()} if self.diagnostic_summary is not None else {}), } @classmethod @@ -484,9 +582,42 @@ def from_dict(cls, data: Mapping[str, object]) -> KernelCheckResult: reason=_optional_str(data, "reason"), returncode=_optional_int(data, "returncode"), diagnostics_truncated=_optional_bool(data, "diagnostics_truncated", default=False), + diagnostic_summary=_diagnostic_summary(data), ) +def aggregate_diagnostic_summaries( + results: tuple[KernelCheckResult, ...], +) -> DiagnosticSummary | None: + """Never present a partial sum as a check-level exact total.""" + summaries = [result.diagnostic_summary for result in results] + if not summaries or not any(summary is not None for summary in summaries): + return None + all_present = all(summary is not None for summary in summaries) + known = [summary for summary in summaries if summary is not None] + healthy = all(result.state is ExecutionState.RAN for result in results) + exact = healthy and all_present and all(s.count_meaning is CountMeaning.EXACT for s in known) + records = (sum(s.emitted_records for s in known) + if all_present and all(s.emitted_records is not None for s in known) else None) + limits = {s.requested_limit for s in known} + + def status(key: str) -> EvidenceStatus: + values = [getattr(s, key) for s in known] + if EvidenceStatus.INCOMPLETE in values: + return EvidenceStatus.INCOMPLETE + if healthy and all_present and all(value is EvidenceStatus.COMPLETE for value in values): + return EvidenceStatus.COMPLETE + return EvidenceStatus.UNKNOWN + + return DiagnosticSummary( + reported_count=sum(s.reported_count for s in known) if exact else None, + count_meaning=CountMeaning.EXACT if exact else CountMeaning.UNKNOWN, + emitted_records=records, + requested_limit=next(iter(limits)) if all_present and len(limits) == 1 else None, + details=status("details"), native_analysis=status("native_analysis"), + ) + + @dataclass(frozen=True) class CheckResult: sanitizer: str @@ -500,9 +631,9 @@ class CheckResult: kernel_results: tuple[KernelCheckResult, ...] = () coverage: tuple[ObjectCoverage, ...] = () backend: tuple[tuple[str, str], ...] = () - # True when any scanned kernel hit the backend's diagnostic cap, so this - # check's aggregated ``findings`` under-counts the hazards present (#480). + # Historical OR of per-entry >= marker flags; false is not completeness. diagnostics_truncated: bool = False + diagnostic_summary: DiagnosticSummary | None = None def __post_init__(self) -> None: if not self.sanitizer: @@ -532,6 +663,14 @@ def __post_init__(self) -> None: result.diagnostics_truncated for result in self.kernel_results ): raise ValueError("a check cannot report a complete count over a truncated kernel") + if self.diagnostic_summary is not None: + if not isinstance(self.diagnostic_summary, DiagnosticSummary): + raise TypeError("diagnostic_summary must be DiagnosticSummary or null") + if (self.diagnostics_truncated + and self.diagnostic_summary.count_meaning is CountMeaning.EXACT): + raise ValueError("historical lower-bound marker contradicts exact count") + if self.diagnostic_summary != aggregate_diagnostic_summaries(self.kernel_results): + raise ValueError("check diagnostic_summary contradicts per-entry evidence") def to_dict(self) -> dict[str, object]: return { @@ -545,6 +684,7 @@ def to_dict(self) -> dict[str, object]: "coverage": [item.to_dict() for item in self.coverage], "backend": dict(self.backend), "diagnostics_truncated": self.diagnostics_truncated, + **({"diagnostic_summary": self.diagnostic_summary.to_dict()} if self.diagnostic_summary is not None else {}), } @classmethod @@ -576,6 +716,7 @@ def from_dict(cls, data: Mapping[str, object]) -> CheckResult: sorted((key, value) for key, value in backend.items() if isinstance(value, str)) ), diagnostics_truncated=_optional_bool(data, "diagnostics_truncated", default=False), + diagnostic_summary=_diagnostic_summary(data), ) @@ -609,6 +750,40 @@ def __post_init__(self) -> None: names = [check.sanitizer for check in self.checks] if len(names) != len(set(names)): raise ValueError("report contains duplicate sanitizer results") + for check in self.checks: + summary = check.diagnostic_summary + if summary is None: + continue + # New affirmative aggregate claims require the full selected scope. + # Whole-object scans intentionally collapse aliases of one object. + def matches(left: KernelIdentity, right: KernelIdentity) -> bool: + return left == right or ( + left.code_object_scan and right.code_object_scan + and left.target == right.target + and left.code_object_sha256 == right.code_object_sha256 + and left.code_object_index == right.code_object_index + ) + + if (summary.count_meaning is CountMeaning.EXACT + or EvidenceStatus.COMPLETE in (summary.details, summary.native_analysis)): + if not self.worklist.kernels or not all( + any(matches(result.identity, item.identity) for result in check.kernel_results) + for item in self.worklist.kernels + ): + raise ValueError("aggregate diagnostic evidence is missing selected results") + if any( + not any(matches(result.identity, item.identity) for item in self.worklist.kernels) + for result in check.kernel_results + ): + raise ValueError("aggregate diagnostic evidence includes unselected results") + keys = [ + ("object", result.identity.target, result.identity.code_object_sha256, + result.identity.code_object_index) + if result.identity.code_object_scan else ("entry", result.identity.stable_key) + for result in check.kernel_results + ] + if len(keys) != len(set(keys)): + raise ValueError("aggregate diagnostic evidence contains duplicate results") verdict = ( Verdict.NOT_CHECKED if not self.checks diff --git a/src/aorta/instrumentation/rocjitsu_sanitizers/recipe.py b/src/aorta/instrumentation/rocjitsu_sanitizers/recipe.py index 6a9f49763..a2c37c731 100644 --- a/src/aorta/instrumentation/rocjitsu_sanitizers/recipe.py +++ b/src/aorta/instrumentation/rocjitsu_sanitizers/recipe.py @@ -261,11 +261,11 @@ def _optional_timeout_seconds(block: Mapping[str, object]) -> float | None: def _optional_waitcheck_max_diagnostics(block: Mapping[str, object]) -> int | None: """Parse an optional positive ``waitcheck_max_diagnostics`` from a block. - Absent keeps the backend's own cap (``rj_waitcheck --max-diagnostics``, 32 - per code object), which reports a capped scan as ``diagnostics_truncated``. + Absent keeps the backend's own default. The requested per-invocation detail + limit is not proof of count or native analysis completeness. A bool, non-integer, or non-positive value is rejected at load so a malformed knob cannot silently reinstate the default cap on a lane that - asked for a complete count. + requested a different diagnostic detail limit. """ if "waitcheck_max_diagnostics" not in block: diff --git a/src/aorta/instrumentation/rocjitsu_sanitizers/waitcheck.py b/src/aorta/instrumentation/rocjitsu_sanitizers/waitcheck.py index 6158147cf..5091a4399 100644 --- a/src/aorta/instrumentation/rocjitsu_sanitizers/waitcheck.py +++ b/src/aorta/instrumentation/rocjitsu_sanitizers/waitcheck.py @@ -9,13 +9,16 @@ import shutil from collections.abc import Mapping, Sequence from concurrent.futures import ThreadPoolExecutor -from dataclasses import dataclass +from dataclasses import dataclass, replace from pathlib import Path from typing import Protocol from .execution import ProcessResult, run_argv from .models import ( CheckResult, + CountMeaning, + DiagnosticSummary, + EvidenceStatus, ExecutionState, Finding, FindingSeverity, @@ -23,41 +26,39 @@ KernelIdentity, KernelWorklist, Verdict, + aggregate_diagnostic_summaries, ) WAITCHECK_CLEAN_EXIT = 0 WAITCHECK_HAZARD_EXIT = 4 WAITCHECK_JSON_SCHEMA = "rj-waitcheck-diagnostic-v1" _HEADER = re.compile( - r":(?Pgfx[0-9a-f]+)\[(?P[0-9]+)\]:" + r"^(?:rj_waitcheck: )?(?P.+):(?Pgfx[0-9a-f]+)\[(?P[0-9]+)\]:" # An exact-entry run prefixes the summary with a kernel marker such as # ``kernel=.text+0x120:``; accept and capture it so a real exact-entry run is # not misread as a missing analysis summary. r"(?:\s*kernel=\.text\+0x(?P[0-9a-f]+):)?" - r"\s+instructions=(?P[0-9]+).*" - # rj_waitcheck stops collecting at --max-diagnostics and marks the capped - # summary by prefixing the count with ">=" (``diagnostics=>=32``). Capture - # that marker: matching it and dropping it made a partial hazard count - # indistinguishable from a complete one (#480). - r"diagnostics=(?P>=)?(?P[0-9]+)" + r"\s+instructions=(?P>=)?(?P[0-9]+) " + r"memory-events=(?P>=)?(?P[0-9]+) " + r"diagnostics=(?P>=)?(?P[0-9]+)$" ) _DIAGNOSTIC = re.compile(r"\bmissing\s+s_wait|\bhazard\b", re.IGNORECASE) _CONTEXT = re.compile(r"^(producer|consumer)\b", re.IGNORECASE) +_OMITTED = re.compile(r"^omitted (?Pat least )?(?P[0-9]+) diagnostic\(s\) after limit$") +_STOPPED = "analysis stopped early; additional diagnostics may exist" @dataclass(frozen=True) class ParsedWaitcheckOutput: - """One Waitcheck scan's findings and whether the backend capped them. + """Unique findings, historical >= marker, and optional native evidence. - ``diagnostics_truncated`` is true when the analysis summary reported - ``diagnostics=>=N``: the scan hit ``--max-diagnostics`` (32 by default) and - ``findings`` is a partial view of the hazards in the code object. Raise the - cap -- ``sanitizer_plan.policy.waitcheck_max_diagnostics`` on a recipe -- to - get a complete count. + False truncation and successful parsing never prove exhaustive analysis. + The default preserves existing two-argument constructors. """ findings: tuple[Finding, ...] diagnostics_truncated: bool + diagnostic_summary: DiagnosticSummary | None = None class ProcessExecutor(Protocol): @@ -224,22 +225,32 @@ def parse_waitcheck_text( *, expected: KernelIdentity, ) -> ParsedWaitcheckOutput: - """Parse exact-entry CLI output until structured single-entry output exists. + """Preserve counts and raw records before deduplication, failing on conflicts. - Returns the findings together with the analysis summary's ``>=`` cap marker, - so a capped scan is not read as a complete one (``ParsedWaitcheckOutput``). + Bare counts have unknown semantics for an unidentified binary. Explicit + omission/stop messages follow rocm-systems d03c6fe; see the reporting note. """ findings: list[Finding] = [] - header_seen = False + summary = None truncated = False + signal = None + omitted = None + early = False expected_index = expected.code_object_index or 0 for raw_line in output.splitlines(): line = raw_line.strip() header = _HEADER.search(line) if header is not None: - header_seen = True - truncated = truncated or header.group("truncated") is not None + if summary is not None: + raise ValueError("Waitcheck output contains multiple/conflicting analysis summaries") + summary = header + truncated = header.group("truncated") is not None + bounds = (bool(header.group("instruction_bound")), bool(header.group("event_bound"))) + if any(bounds) and not (all(bounds) and truncated): + raise ValueError("Waitcheck instruction/event bounds contradict diagnostic bound") + if header.group("object") != expected.code_object: + raise ValueError("Waitcheck object does not match requested code object") if header.group("target") != expected.target: raise ValueError( f"Waitcheck target {header.group('target')!r} does not match " @@ -261,8 +272,24 @@ def parse_waitcheck_text( f"requested 0x{expected.entry_offset:x}" ) continue + if "instructions=" in line or "diagnostics=" in line: + raise ValueError("Waitcheck output contains a malformed analysis summary") + omission = _OMITTED.fullmatch(line) + if omission is not None or line == _STOPPED: + if summary is None or signal is not None: + raise ValueError("Waitcheck completion signal missing summary or conflicting") + signal = line + early = line == _STOPPED or omission.group("bound") is not None + omitted = int(omission.group("count")) if omission is not None else None + if early != truncated: + raise ValueError("Waitcheck completion signal contradicts count bound") + if any(bool(summary.group(key)) != early for key in ("instruction_bound", "event_bound")): + raise ValueError("Waitcheck explicit signal contradicts instruction/event bounds") + continue context = _CONTEXT.match(line) if context is not None: + if not findings or signal is not None: + raise ValueError("Waitcheck context without an active diagnostic") if findings: previous = findings[-1] context_index = ( @@ -287,6 +314,11 @@ def parse_waitcheck_text( ) continue if _DIAGNOSTIC.search(line): + prefix = re.escape(f"{expected.code_object}:{expected.target}[{expected_index}]:") + if not re.match(prefix + r"\.text\+0x[0-9a-f]+: ", line): + raise ValueError("Waitcheck diagnostic object/target/image prefix does not match requested identity") + if summary is None or signal is not None: + raise ValueError("Waitcheck diagnostic outside summary/detail sequence") findings.append( Finding( sanitizer="waitcheck", @@ -301,12 +333,40 @@ def parse_waitcheck_text( entry_offset=expected.entry_offset, ) ) - if not header_seen: + continue + if line: + if summary is None: + raise ValueError("Waitcheck output did not contain an analysis summary; unrecognized line: " + line) + raise ValueError(f"Waitcheck output contains an unrecognized line: {line}") + if summary is None: raise ValueError("Waitcheck output did not contain an analysis summary") + count = int(summary.group("diagnostics")) + records = len(findings) + if records > count: + raise ValueError("Waitcheck summary count is smaller than emitted diagnostic records") + if omitted is not None and (omitted < 1 or count - records != omitted): + raise ValueError("Waitcheck omitted count contradicts summary and emitted records") + if signal == _STOPPED and records != count: + raise ValueError("Waitcheck early-stop signal does not account for omitted records") + meaning = CountMeaning.LOWER_BOUND if truncated else CountMeaning.UNKNOWN + details = EvidenceStatus.UNKNOWN + native = EvidenceStatus.UNKNOWN + if signal is not None: + meaning = CountMeaning.LOWER_BOUND if early else CountMeaning.EXACT + native = EvidenceStatus.INCOMPLETE if early else EvidenceStatus.COMPLETE + details = EvidenceStatus.INCOMPLETE if omitted is not None else EvidenceStatus.COMPLETE + # The early-stop sentence accounts for observed details only, not the + # diagnostics that unvisited code could have produced. + if early and omitted is None: + details = EvidenceStatus.UNKNOWN deduplicated = {finding.dedupe_key: finding for finding in findings} return ParsedWaitcheckOutput( findings=tuple(deduplicated[key] for key in sorted(deduplicated, key=repr)), diagnostics_truncated=truncated, + diagnostic_summary=DiagnosticSummary( + reported_count=count, count_meaning=meaning, emitted_records=records, + details=details, native_analysis=native, + ), ) @@ -353,12 +413,16 @@ def waitcheck_argv( return tuple(argv) -def _not_checked(identity: KernelIdentity, reason: str) -> KernelCheckResult: +def _not_checked( + identity: KernelIdentity, reason: str, max_diagnostics: int | None = None, +) -> KernelCheckResult: return KernelCheckResult( identity=identity, state=ExecutionState.NOT_CHECKED, verdict=Verdict.NOT_CHECKED, reason=reason, + diagnostic_summary=(DiagnosticSummary(requested_limit=max_diagnostics) + if max_diagnostics is not None else None), ) @@ -371,13 +435,14 @@ def _run_one( execute: ProcessExecutor, max_diagnostics: int | None = None, ) -> KernelCheckResult: + requested = DiagnosticSummary(requested_limit=max_diagnostics) if max_diagnostics is not None else None if not identity.exact and not identity.code_object_scan: - return _not_checked(identity, "code_object_identity_required") + return _not_checked(identity, "code_object_identity_required", max_diagnostics) artifact = Path(str(identity.code_object)) if not artifact.is_file(): - return _not_checked(identity, "code_object_not_found") + return _not_checked(identity, "code_object_not_found", max_diagnostics) if artifact.stat().st_size == 0: - return _not_checked(identity, "code_object_empty") + return _not_checked(identity, "code_object_empty", max_diagnostics) digest = _sha256_file(artifact) if digest != identity.code_object_sha256: return KernelCheckResult( @@ -385,6 +450,7 @@ def _run_one( state=ExecutionState.ERROR, verdict=Verdict.ERROR, reason="code_object_digest_mismatch", + diagnostic_summary=requested, ) process = execute( waitcheck_argv(binary, identity, max_diagnostics=max_diagnostics), @@ -400,6 +466,7 @@ def _run_one( state=ExecutionState.TIMED_OUT, verdict=Verdict.ERROR, reason="waitcheck_timeout", + diagnostic_summary=requested, ) if process.launch_error is not None: return KernelCheckResult( @@ -407,6 +474,7 @@ def _run_one( state=ExecutionState.ERROR, verdict=Verdict.ERROR, reason=f"waitcheck_launch_error: {process.launch_error}", + diagnostic_summary=requested, ) # Parse defensively so an unexpected exit code can still be reported with the # backend's own message. When rj_waitcheck refuses an input it explains itself on @@ -418,6 +486,7 @@ def _run_one( parse_error: Exception | None = None findings: tuple[Finding, ...] = () diagnostics_truncated = False + diagnostic_summary = requested try: parsed = parse_waitcheck_text( f"{process.stdout}\n{process.stderr}", @@ -425,6 +494,8 @@ def _run_one( ) findings = parsed.findings diagnostics_truncated = parsed.diagnostics_truncated + if parsed.diagnostic_summary is not None: + diagnostic_summary = replace(parsed.diagnostic_summary, requested_limit=max_diagnostics) except Exception as exc: # noqa: BLE001 - see below # Deliberately broad. The point of parsing here is that an unexpected exit code # reaches the branch below with the backend's own message; a parser raising @@ -445,6 +516,7 @@ def _run_one( returncode=process.returncode, findings=findings, diagnostics_truncated=diagnostics_truncated, + diagnostic_summary=diagnostic_summary, ) if parse_error is not None: # An expected exit code with unparsable output is a real identity/format @@ -455,6 +527,7 @@ def _run_one( verdict=Verdict.ERROR, reason=f"waitcheck_diagnostics_error: {parse_error}", returncode=process.returncode, + diagnostic_summary=requested, ) if process.returncode == WAITCHECK_HAZARD_EXIT and not findings: return KernelCheckResult( @@ -464,6 +537,16 @@ def _run_one( reason="waitcheck_hazard_exit_without_structured_diagnostics", returncode=process.returncode, diagnostics_truncated=diagnostics_truncated, + diagnostic_summary=diagnostic_summary, + ) + if process.returncode == WAITCHECK_CLEAN_EXIT and ( + findings or (diagnostic_summary is not None and diagnostic_summary.reported_count) + ): + return KernelCheckResult( + identity=identity, state=ExecutionState.ERROR, verdict=Verdict.ERROR, + reason="waitcheck_clean_exit_with_diagnostics", returncode=process.returncode, + findings=findings, diagnostics_truncated=diagnostics_truncated, + diagnostic_summary=diagnostic_summary, ) return KernelCheckResult( identity=identity, @@ -472,6 +555,7 @@ def _run_one( returncode=process.returncode, findings=findings, diagnostics_truncated=diagnostics_truncated, + diagnostic_summary=diagnostic_summary, ) @@ -486,10 +570,9 @@ def run_waitcheck( ) -> CheckResult: """Run exact Waitcheck for every worklist entry, failing closed on gaps. - ``max_diagnostics`` overrides the backend's per-code-object diagnostic cap - (``rj_waitcheck --max-diagnostics``, 32 by default). Leaving it unset keeps - the backend default; either way a capped scan is reported as - ``diagnostics_truncated`` rather than as a complete hazard count. + ``max_diagnostics`` requests a per-invocation detail limit. Leaving it unset + keeps the backend default. Count/completion semantics depend on the format; + ``diagnostics_truncated`` retains only the historical >= marker meaning. """ # Validate at the public boundary so missing backends and empty/rejected @@ -503,7 +586,7 @@ def run_waitcheck( verdict=Verdict.NOT_CHECKED, reason="rj_waitcheck_not_found", kernel_results=tuple( - _not_checked(item.identity, "rj_waitcheck_not_found") for item in worklist.kernels + _not_checked(item.identity, "rj_waitcheck_not_found", max_diagnostics) for item in worklist.kernels ), ) if resolved.stat().st_size == 0: @@ -513,7 +596,7 @@ def run_waitcheck( verdict=Verdict.NOT_CHECKED, reason="rj_waitcheck_empty", kernel_results=tuple( - _not_checked(item.identity, "rj_waitcheck_empty") for item in worklist.kernels + _not_checked(item.identity, "rj_waitcheck_empty", max_diagnostics) for item in worklist.kernels ), ) backend = ( @@ -565,8 +648,7 @@ def _scan(task: tuple[int, KernelIdentity]) -> KernelCheckResult: with ThreadPoolExecutor(max_workers=max_workers) as pool: kernel_results = tuple(pool.map(_scan, scan_tasks)) findings = tuple(finding for result in kernel_results for finding in result.findings) - # One capped scan makes the whole check's finding count a floor, not a total, - # so the flag rolls up alongside the findings it qualifies. + # Preserve the historical marker independently from new native evidence. truncated = any(result.diagnostics_truncated for result in kernel_results) if not kernel_results: return CheckResult( @@ -591,6 +673,7 @@ def _scan(task: tuple[int, KernelIdentity]) -> KernelCheckResult: kernel_results=kernel_results, backend=backend, diagnostics_truncated=truncated, + diagnostic_summary=aggregate_diagnostic_summaries(kernel_results), ) return CheckResult( sanitizer="waitcheck", @@ -600,4 +683,5 @@ def _scan(task: tuple[int, KernelIdentity]) -> KernelCheckResult: kernel_results=kernel_results, backend=backend, diagnostics_truncated=truncated, + diagnostic_summary=aggregate_diagnostic_summaries(kernel_results), ) diff --git a/tests/instrumentation/rocjitsu_sanitizers/fixtures/README.md b/tests/instrumentation/rocjitsu_sanitizers/fixtures/README.md new file mode 100644 index 000000000..d36e28c0f --- /dev/null +++ b/tests/instrumentation/rocjitsu_sanitizers/fixtures/README.md @@ -0,0 +1,16 @@ +# Waitcheck reporting fixtures + +Minimal text derivatives of the retained 11 November 2026 cap64/cap1 control +logs. The object path alone is replaced by `/fixtures/multiple_missing_waits`; +the summary numbers, gfx1250/image0/entry0, diagnostic offsets/instructions and +producer/consumer text are retained. No binary or private path is included. +Whitespace is normalized to LF. Counts/observations were previously established +by the saved `inspect_partial.py` analyzer. The updated saved reader additionally +recognizes the cap1 terminal `omitted 2 diagnostic(s) after limit`; retain it. +The old reader's count gate had failed before considering this line. A synthetic +test separately removes the signal to verify that bare mismatches stay unknown. +Fixtures are parser data, never run. + +The tests substitute the portable object path with an inert temporary file when +checking recipe/backend integration. Native4 is supplied by a fake executor. +These tests do not establish original binary provenance or native qualification. diff --git a/tests/instrumentation/rocjitsu_sanitizers/fixtures/waitcheck-cap1.txt b/tests/instrumentation/rocjitsu_sanitizers/fixtures/waitcheck-cap1.txt new file mode 100644 index 000000000..52462e64a --- /dev/null +++ b/tests/instrumentation/rocjitsu_sanitizers/fixtures/waitcheck-cap1.txt @@ -0,0 +1,5 @@ +rj_waitcheck: /fixtures/multiple_missing_waits:gfx1250[0]:kernel=.text+0x0: instructions=10 memory-events=6 diagnostics=3 +/fixtures/multiple_missing_waits:gfx1250[0]:.text+0x8: missing s_wait_kmcnt <= 0 before use of s4 + producer .text+0x0: s_load_b32 s4, s[0:1], 0x0 + consumer .text+0x8: s_mov_b32 s5, s4 +omitted 2 diagnostic(s) after limit diff --git a/tests/instrumentation/rocjitsu_sanitizers/fixtures/waitcheck-cap64.txt b/tests/instrumentation/rocjitsu_sanitizers/fixtures/waitcheck-cap64.txt new file mode 100644 index 000000000..5087a24ec --- /dev/null +++ b/tests/instrumentation/rocjitsu_sanitizers/fixtures/waitcheck-cap64.txt @@ -0,0 +1,10 @@ +rj_waitcheck: /fixtures/multiple_missing_waits:gfx1250[0]:kernel=.text+0x0: instructions=10 memory-events=6 diagnostics=3 +/fixtures/multiple_missing_waits:gfx1250[0]:.text+0x8: missing s_wait_kmcnt <= 0 before use of s4 + producer .text+0x0: s_load_b32 s4, s[0:1], 0x0 + consumer .text+0x8: s_mov_b32 s5, s4 +/fixtures/multiple_missing_waits:gfx1250[0]:.text+0x18: missing s_wait_kmcnt <= 0 before use of s6 + producer .text+0x10: s_load_b32 s6, s[0:1], 0x4 + consumer .text+0x18: s_mov_b32 s7, s6 +/fixtures/multiple_missing_waits:gfx1250[0]:.text+0x28: missing s_wait_kmcnt <= 0 before use of s8 + producer .text+0x20: s_load_b32 s8, s[0:1], 0x8 + consumer .text+0x28: s_mov_b32 s9, s8 diff --git a/tests/instrumentation/rocjitsu_sanitizers/test_diagnostic_summary.py b/tests/instrumentation/rocjitsu_sanitizers/test_diagnostic_summary.py new file mode 100644 index 000000000..cf81be68c --- /dev/null +++ b/tests/instrumentation/rocjitsu_sanitizers/test_diagnostic_summary.py @@ -0,0 +1,185 @@ +"""Optional diagnostic evidence validation and conservative aggregation.""" +from dataclasses import replace + +import pytest + +from aorta.instrumentation.rocjitsu_sanitizers.models import ( + CheckResult, CountMeaning, DiagnosticSummary, EvidenceStatus, ExecutionState, + KernelCheckResult, KernelIdentity, aggregate_diagnostic_summaries, Verdict, + KernelObservation, KernelWorklist, SanitizerReport, SelectionRequirement, +) + + +def entry(summary=None, state=ExecutionState.RAN): + return KernelCheckResult( + identity=KernelIdentity(name="test", target="gfx1250"), state=state, + verdict=Verdict.WARN if state is ExecutionState.RAN else Verdict.ERROR, + diagnostic_summary=summary, + ) + + +EXACT = DiagnosticSummary(reported_count=3, count_meaning=CountMeaning.EXACT, + emitted_records=1, requested_limit=1, + details=EvidenceStatus.INCOMPLETE, native_analysis=EvidenceStatus.COMPLETE) + + +def test_old_reports_and_constructors_have_no_affirmative_evidence(): + old = entry() + assert "diagnostic_summary" not in old.to_dict() + assert KernelCheckResult.from_dict(old.to_dict()) == old + check = CheckResult("waitcheck", ExecutionState.RAN, Verdict.WARN, kernel_results=(old,)) + assert "diagnostic_summary" not in check.to_dict() + assert CheckResult.from_dict(check.to_dict()) == check + assert DiagnosticSummary.from_dict({}) == DiagnosticSummary() + + +def test_optional_metadata_round_trip(): + result = entry(EXACT) + summary = aggregate_diagnostic_summaries((result,)) + check = CheckResult("waitcheck", ExecutionState.RAN, Verdict.WARN, + kernel_results=(result,), diagnostic_summary=summary) + assert CheckResult.from_dict(check.to_dict()) == check + + +@pytest.mark.parametrize("field,value", [ + ("reported_count", True), ("reported_count", -1), ("reported_count", 1.0), + ("emitted_records", "1"), ("emitted_records", -1), ("requested_limit", 0), + ("requested_limit", False), ("count_meaning", "total"), ("details", "yes"), + ("native_analysis", "success"), ("unknown", 1), +]) +def test_malformed_metadata_is_rejected(field, value): + with pytest.raises((ValueError, TypeError)): + DiagnosticSummary.from_dict({field: value}) + + +@pytest.mark.parametrize("changes", [ + {"reported_count": None}, {"reported_count": 0}, + {"details": EvidenceStatus.COMPLETE}, + {"count_meaning": CountMeaning.LOWER_BOUND}, + {"native_analysis": EvidenceStatus.INCOMPLETE}, +]) +def test_contradictory_metadata_is_rejected(changes): + with pytest.raises(ValueError): + replace(EXACT, **changes) + + +@pytest.mark.parametrize("other", [ + entry(), entry(DiagnosticSummary(reported_count=3, emitted_records=1)), + entry(DiagnosticSummary(reported_count=3, count_meaning=CountMeaning.LOWER_BOUND)), + entry(EXACT, ExecutionState.ERROR), +]) +def test_mixed_unknown_missing_lower_bound_and_error_do_not_fabricate_totals(other): + summary = aggregate_diagnostic_summaries((entry(EXACT), other)) + assert summary.reported_count is None + assert summary.count_meaning is CountMeaning.UNKNOWN + assert summary.native_analysis is not EvidenceStatus.COMPLETE + + +def test_exact_aggregate_and_explicit_early_stop(): + summary = aggregate_diagnostic_summaries((entry(EXACT), entry(EXACT))) + assert (summary.reported_count, summary.emitted_records, summary.requested_limit) == (6, 2, 1) + assert summary.native_analysis is EvidenceStatus.COMPLETE + assert summary.details is EvidenceStatus.INCOMPLETE + early = replace(EXACT, count_meaning=CountMeaning.LOWER_BOUND, native_analysis=EvidenceStatus.INCOMPLETE) + assert aggregate_diagnostic_summaries((entry(EXACT), entry(early))).native_analysis is EvidenceStatus.INCOMPLETE + + +def test_invalid_aggregate_claim_is_rejected(): + with pytest.raises(ValueError, match="per-entry"): + CheckResult("waitcheck", ExecutionState.RAN, Verdict.WARN, + kernel_results=(entry(),), diagnostic_summary=EXACT) + + +def test_exact_report_requires_results_for_the_full_worklist(): + result = entry(EXACT) + check = CheckResult("waitcheck", ExecutionState.RAN, Verdict.WARN, + kernel_results=(result,), diagnostic_summary=EXACT) + worklist = KernelWorklist( + requirement=SelectionRequirement.TOP_TIME, top_n=2, + kernels=tuple(KernelObservation(identity=KernelIdentity(name=name, target="gfx1250"), + sources=("synthetic",)) for name in ("test", "missing")), + ) + with pytest.raises(ValueError, match="missing selected"): + SanitizerReport(target="gfx1250", worklist=worklist, checks=(check,)) + + +def test_complete_details_and_native_analysis_stay_separate(): + summary = DiagnosticSummary(reported_count=1, count_meaning=CountMeaning.EXACT, + emitted_records=1, details=EvidenceStatus.COMPLETE) + aggregate = aggregate_diagnostic_summaries((entry(summary),)) + assert aggregate.details is EvidenceStatus.COMPLETE + assert aggregate.native_analysis is EvidenceStatus.UNKNOWN + + +def test_exact_report_cannot_include_unselected_results(): + first = entry(EXACT) + foreign = replace(first, identity=KernelIdentity(name="foreign", target="gfx1250")) + results = (first, foreign) + check = CheckResult("waitcheck", ExecutionState.RAN, Verdict.WARN, + kernel_results=results, diagnostic_summary=aggregate_diagnostic_summaries(results)) + worklist = KernelWorklist(requirement=SelectionRequirement.TOP_TIME, top_n=1, + kernels=(KernelObservation(identity=first.identity, sources=("test",)),)) + with pytest.raises(ValueError, match="unselected"): + SanitizerReport(target="gfx1250", worklist=worklist, checks=(check,)) + + +def test_exact_aggregate_rejects_duplicate_whole_object_aliases(): + identity = KernelIdentity(name="first", target="gfx1250", code_object="/fixtures/object.co", + code_object_sha256="a" * 64, code_object_index=0) + first = replace(entry(EXACT), identity=identity) + alias = replace(first, identity=replace(identity, name="alias")) + results = (first, alias) + check = CheckResult("waitcheck", ExecutionState.RAN, Verdict.WARN, kernel_results=results, + diagnostic_summary=aggregate_diagnostic_summaries(results)) + worklist = KernelWorklist(requirement=SelectionRequirement.TOP_TIME, top_n=1, + kernels=(KernelObservation(identity=identity, sources=("test",)),)) + with pytest.raises(ValueError, match="duplicate results"): + SanitizerReport(target="gfx1250", worklist=worklist, checks=(check,)) + + +def test_one_whole_object_result_covers_selected_aliases_without_double_counting(): + identity = KernelIdentity(name="first", target="gfx1250", code_object="/fixtures/object.co", + code_object_sha256="a" * 64, code_object_index=0) + result = replace(entry(EXACT), identity=identity) + check = CheckResult("waitcheck", ExecutionState.RAN, Verdict.WARN, kernel_results=(result,), + diagnostic_summary=aggregate_diagnostic_summaries((result,))) + worklist = KernelWorklist(requirement=SelectionRequirement.TOP_TIME, top_n=2, + kernels=tuple(KernelObservation(identity=item, sources=("test",)) + for item in (identity, replace(identity, name="alias")))) + report = SanitizerReport(target="gfx1250", worklist=worklist, checks=(check,)) + assert report.checks[0].diagnostic_summary.reported_count == 3 + + +@pytest.mark.parametrize('entry_flag', ['omitted', False]) +def test_check_rejects_exact_summary_with_aggregate_lower_bound(entry_flag): + result = entry(EXACT) + data = CheckResult('waitcheck', ExecutionState.RAN, Verdict.WARN, + kernel_results=(result,), diagnostic_summary=EXACT).to_dict() + data['diagnostics_truncated'] = True + if entry_flag == 'omitted': + data['kernel_results'][0].pop('diagnostics_truncated') + with pytest.raises(ValueError, match='lower-bound marker contradicts exact count'): + CheckResult.from_dict(data) + + +@pytest.mark.parametrize('summary', [None, DiagnosticSummary(), + DiagnosticSummary(reported_count=3, count_meaning=CountMeaning.LOWER_BOUND)]) +def test_check_keeps_legacy_truncation_with_nonexact_or_absent_summary(summary): + result = entry(summary) + aggregate = aggregate_diagnostic_summaries((result,)) + check = CheckResult('waitcheck', ExecutionState.RAN, Verdict.WARN, + kernel_results=(result,), diagnostic_summary=aggregate, + diagnostics_truncated=True) + assert CheckResult.from_dict(check.to_dict()) == check + + +def test_multi_entry_check_can_exceed_the_per_invocation_limit(): + first = entry(EXACT) + second = replace(first, identity=KernelIdentity(name='second', target='gfx1250')) + results = (first, second) + aggregate = aggregate_diagnostic_summaries(results) + check = CheckResult('waitcheck', ExecutionState.RAN, Verdict.WARN, + kernel_results=results, diagnostic_summary=aggregate) + assert check.diagnostic_summary.emitted_records == 2 + assert check.diagnostic_summary.requested_limit == 1 + assert CheckResult.from_dict(check.to_dict()) == check diff --git a/tests/instrumentation/rocjitsu_sanitizers/test_waitcheck.py b/tests/instrumentation/rocjitsu_sanitizers/test_waitcheck.py index 5bc93eeb4..af151fce2 100644 --- a/tests/instrumentation/rocjitsu_sanitizers/test_waitcheck.py +++ b/tests/instrumentation/rocjitsu_sanitizers/test_waitcheck.py @@ -10,6 +10,8 @@ import pytest from aorta.instrumentation.rocjitsu_sanitizers import ( + CountMeaning, + EvidenceStatus, ExecutionState, KernelIdentity, KernelObservation, @@ -99,7 +101,7 @@ def execute( assert "--code-object-index" in argv assert argv[argv.index("--kernel-entry") + 1] == "0x120" assert "--diagnostics-jsonl" not in argv - # An uncapped summary is a complete count, and nothing overrides the + # A bare summary does not prove exhaustive analysis. Nothing overrides the # backend's own diagnostic cap unless a caller asks for it. assert "--max-diagnostics" not in argv assert result.diagnostics_truncated is False @@ -152,8 +154,7 @@ def execute( def test_waitcheck_raises_the_backend_cap_when_asked(tmp_path: Path) -> None: - # The cap is what truncates, so a lane that needs a complete count must be able - # to raise it without editing code. + # The requested detail cap reaches argv and does not establish completeness. binary = tmp_path / "rj_waitcheck" binary.write_text("binary") artifact = tmp_path / "library.so" @@ -169,7 +170,7 @@ def execute( captured.append(tuple(argv)) output = ( f"{artifact}:gfx950[3]:kernel=.text+0x120: " - "instructions=43023 memory-events=5198 diagnostics=45" + "instructions=43023 memory-events=5198 diagnostics=0" ) return ProcessResult(tuple(argv), 0, output, "") @@ -184,6 +185,8 @@ def execute( argv = captured[0] assert argv[argv.index("--max-diagnostics") + 1] == "100000" assert result.diagnostics_truncated is False + assert result.verdict is Verdict.PASS + assert result.kernel_results[0].diagnostic_summary.requested_limit == 100000 @pytest.mark.parametrize("bad_cap", [0, -1, True, 1.5, "1"]) @@ -544,6 +547,99 @@ def test_parse_waitcheck_text_accepts_exact_entry_marker() -> None: assert parsed.diagnostics_truncated is False +@pytest.mark.parametrize("cap,records", [(64, 3), (1, 1)]) +def test_saved_control_derivatives_preserve_count_details_and_signals(tmp_path, cap, records): + artifact = tmp_path / "multiple_missing_waits" + artifact.write_bytes(b"inert object") + binary = tmp_path / "rj_waitcheck" + binary.write_bytes(b"inert backend") + identity = KernelIdentity( + name="waitcheck_multiple_missing", target="gfx1250", code_object=str(artifact), + code_object_sha256=hashlib.sha256(artifact.read_bytes()).hexdigest(), + code_object_index=0, entry_offset=0, + ) + text = (Path(__file__).parent / "fixtures" / f"waitcheck-cap{cap}.txt").read_text() + text = text.replace("/fixtures/multiple_missing_waits", str(artifact)) + calls = [] + + def execute(argv, **kwargs): + calls.append(argv) + return ProcessResult(tuple(argv), 4, text, "") + + check = run_waitcheck(_worklist(identity), output_dir=tmp_path / "out", binary=binary, + max_diagnostics=cap, execute=execute) + entry = check.kernel_results[0] + assert check.verdict is entry.verdict is Verdict.WARN + assert entry.returncode == 4 + assert check.state is ExecutionState.RAN + assert not check.diagnostics_truncated + assert len(entry.findings) == records + summary = entry.diagnostic_summary + assert (summary.reported_count, summary.emitted_records, summary.requested_limit) == (3, records, cap) + if cap == 1: + assert summary.count_meaning is CountMeaning.EXACT + assert summary.details is EvidenceStatus.INCOMPLETE + assert summary.native_analysis is EvidenceStatus.COMPLETE + assert check.diagnostic_summary.reported_count == 3 + else: + assert summary.count_meaning is CountMeaning.UNKNOWN + assert summary.details is summary.native_analysis is EvidenceStatus.UNKNOWN + assert check.diagnostic_summary.reported_count is None + assert len(calls) == 1 # No recovery scan or cap change. + + +def _text_summary(count="1", bound="", signal="", copies=1): + return ( + f"/tmp/library.so:gfx950[3]:kernel=.text+0x120: " + f"instructions={bound}10 memory-events={bound}2 diagnostics={count}\n" + + "/tmp/library.so:gfx950[3]:.text+0x20: missing s_waitcnt lgkmcnt(0)\n" * copies + + signal + ) + + +def test_raw_duplicates_are_not_detail_omissions(): + parsed = parse_waitcheck_text(_text_summary("2", copies=2), expected=_entry_identity()) + assert len(parsed.findings) == 1 + assert parsed.diagnostic_summary.emitted_records == 2 + assert parsed.diagnostic_summary.reported_count == 2 + assert parsed.diagnostic_summary.details is EvidenceStatus.UNKNOWN + assert not parsed.diagnostics_truncated + + +@pytest.mark.parametrize("count,bound,signal,meaning,details,native", [ + ("3", "", "omitted 2 diagnostic(s) after limit", CountMeaning.EXACT, + EvidenceStatus.INCOMPLETE, EvidenceStatus.COMPLETE), + (">=3", ">=", "omitted at least 2 diagnostic(s) after limit", CountMeaning.LOWER_BOUND, + EvidenceStatus.INCOMPLETE, EvidenceStatus.INCOMPLETE), + (">=1", ">=", "analysis stopped early; additional diagnostics may exist", CountMeaning.LOWER_BOUND, + EvidenceStatus.UNKNOWN, EvidenceStatus.INCOMPLETE), + (">=1", "", "", CountMeaning.LOWER_BOUND, EvidenceStatus.UNKNOWN, EvidenceStatus.UNKNOWN), + ("1", "", "", CountMeaning.UNKNOWN, EvidenceStatus.UNKNOWN, EvidenceStatus.UNKNOWN), +]) +def test_explicit_upstream_signals_and_legacy_unknowns(count, bound, signal, meaning, details, native): + parsed = parse_waitcheck_text(_text_summary(count, bound, signal), expected=_entry_identity()) + summary = parsed.diagnostic_summary + assert (summary.count_meaning, summary.details, summary.native_analysis) == (meaning, details, native) + + +@pytest.mark.parametrize("text", [ + _text_summary("0"), + _text_summary() + _text_summary(), + _text_summary().replace("diagnostics=1", "diagnostics=1junk"), + _text_summary("3", signal="omitted 1 diagnostic(s) after limit"), + _text_summary("3", signal="omitted at least 2 diagnostic(s) after limit"), + _text_summary(">=3", ">=", "omitted 2 diagnostic(s) after limit"), + _text_summary(">=3", ">=", "analysis stopped early; additional diagnostics may exist"), + _text_summary("3", signal="omitted 2 diagnostic(s) after limit\nomitted 2 diagnostic(s) after limit"), + _text_summary() + "unknown completion token", + _text_summary().replace("/tmp/library.so", "/foreign/library.so"), + "producer .text+0x0: s_load_b32 s0\n" + _text_summary(), +]) +def test_malformed_or_conflicting_count_evidence_is_rejected(text): + with pytest.raises(ValueError): + parse_waitcheck_text(text, expected=_entry_identity()) + + @pytest.mark.parametrize( ("summary_count", "expected_truncated"), [(">=32", True), ("45", False)], diff --git a/tests/instrumentation/rocjitsu_sanitizers/test_waitcheck_controls_example.py b/tests/instrumentation/rocjitsu_sanitizers/test_waitcheck_controls_example.py new file mode 100644 index 000000000..47bc787ff --- /dev/null +++ b/tests/instrumentation/rocjitsu_sanitizers/test_waitcheck_controls_example.py @@ -0,0 +1,299 @@ +"""Portable example validation with inert inputs and fake AORTA processes only.""" +from __future__ import annotations + +import argparse +import importlib.util +import json +from pathlib import Path +import shutil +import subprocess + +import pytest + +from aorta.instrumentation.rocjitsu_sanitizers.execution import ProcessResult +from aorta.instrumentation.rocjitsu_sanitizers.models import KernelCheckResult, aggregate_diagnostic_summaries +from aorta.instrumentation.rocjitsu_sanitizers.recipe import load_sanitizer_recipe, _resolve_observations +from aorta.instrumentation.rocjitsu_sanitizers.report import build_report, write_report +from aorta.instrumentation.rocjitsu_sanitizers.selection import select_kernels +from aorta.instrumentation.rocjitsu_sanitizers.waitcheck import run_waitcheck + +ROOT = Path(__file__).resolve().parents[3] +SCRIPT = ROOT / "examples/sanitizers/waitcheck/waitcheck_controls.py" +SPEC = importlib.util.spec_from_file_location("waitcheck_controls_example", SCRIPT) +helper = importlib.util.module_from_spec(SPEC) +SPEC.loader.exec_module(helper) + + +@pytest.fixture +def prepared(tmp_path): + missing, corrected = tmp_path / "original-missing", tmp_path / "original-corrected" + missing.write_bytes(b"inert missing control fixture") + corrected.write_bytes(b"inert corrected control fixture") + output = tmp_path / "prepared" + args = argparse.Namespace(output=output, missing_object=missing, corrected_object=corrected, + kernel_name="test_control", build_id="synthetic-no-compiler", + missing_index=2, missing_entry=16, corrected_index=3, corrected_entry=32) + helper.prepare(args) + return output + + +@pytest.fixture +def completed(tmp_path, prepared, monkeypatch): + # Linux platform.platform() may probe `uname -p` when its cache is cold. + # Keep metadata deterministic so the process mock remains AORTA-only. + monkeypatch.setattr(helper.platform, "_platform_cache", None) + monkeypatch.setattr(helper.platform, "platform", lambda: "synthetic-no-probe-platform") + aorta = tmp_path / "aorta" + aorta.write_text("inert fake CLI", encoding="utf-8") + aorta.chmod(0o755) + bundle = tmp_path / "bundle" + (bundle / "bin").mkdir(parents=True) + scanner = bundle / "bin" / "rj_waitcheck" + scanner.write_text("inert fake tool", encoding="utf-8") + scanner.chmod(0o755) + calls, backend_calls = [], [] + + def fake_popen(argv, **kwargs): + calls.append(argv) + assert argv[:4] == [str(aorta), "sweep", "run", "--recipe"] + assert kwargs["shell"] is False + assert kwargs["env"]["ROCJITSU_PREBUILT"] == str(bundle) + recipe = load_sanitizer_recipe(Path(argv[4])) + worklist = select_kernels(_resolve_observations(recipe), requirement=recipe.requirement, top_n=1) + output = Path(argv[6]) + label = Path(argv[4]).stem + + def execute(native_argv, *, timeout_seconds, **unused): + backend_calls.append(native_argv) + assert timeout_seconds == 120 + identity = worklist.kernels[0].identity + count = 1 if label == "missing" else 0 + instructions = 3 if count else 4 + prefix = f"{identity.code_object}:gfx1250[{identity.code_object_index}]" + log = (f"rj_waitcheck: {prefix}:kernel=.text+0x{identity.entry_offset:x}: " + f"instructions={instructions} memory-events=2 diagnostics={count}\n") + if count: + log += (f"{prefix}:.text+0x18: missing s_wait_kmcnt <= 0 before use of s4\n" + " producer .text+0x10: s_load_b32 s4, s[0:1], 0x0\n" + " consumer .text+0x18: s_mov_b32 s5, s4\n") + return ProcessResult(tuple(native_argv), 4 if count else 0, log, "") + + check = run_waitcheck(worklist, output_dir=output / "waitcheck", binary=scanner, + timeout_seconds=120, max_diagnostics=64, execute=execute) + write_report(build_report(target="gfx1250", worklist=worklist, checks=(check,)), + output / "sanitizer_report.json") + + class Process: + def wait(self, timeout): + assert timeout == 180 + return 1 if label == "mismatch" else 0 + return Process() + + monkeypatch.setattr(helper.subprocess, "Popen", fake_popen) + monkeypatch.setattr(helper.importlib.metadata, "version", lambda name: "test-only") + output = tmp_path / "run" + helper.run(argparse.Namespace(prepared=prepared, aorta=aorta, prebuilt=bundle, + tool_sha256=helper.digest(scanner), output=output)) + return output, calls, backend_calls + + +def test_prepare_preserves_inputs_and_explicit_nonzero_scope(prepared): + manifest = helper.check_prepared(prepared) + assert manifest["controls"]["missing"]["index"] == 2 + assert manifest["controls"]["missing"]["entry"] == 16 + assert (prepared.parent / "original-missing").read_bytes() == b"inert missing control fixture" + original = load_sanitizer_recipe(prepared / "missing.json") + assert original.kernel_specs[0].entry_offset == 16 + mismatch = load_sanitizer_recipe(prepared / "mismatch.json") + assert mismatch.kernel_specs[0].code_object_sha256 == "0" * 64 + assert mismatch.kernel_specs[0].entry_offset == 32 + + +def test_three_cases_only_aorta_called_and_mismatch_never_scanned(completed): + output, calls, backend_calls = completed + assert len(calls) == 3 + assert len(backend_calls) == 2 + summary = helper.verify(output) + assert "missing: CLI=0, native=4, state=ran, verdict=warn" in summary + assert "mismatch: CLI=1, native=None, state=error, verdict=error" in summary + assert not (output / "mismatch/waitcheck/waitcheck-0.log").exists() + + +def test_cold_platform_metadata_does_not_escape_aorta_process_mock(completed): + output, calls, backend_calls = completed + manifest = helper.load(output / "run.json") + assert helper.platform._platform_cache is None + assert manifest["platform"] == "synthetic-no-probe-platform" + assert len(calls) == 3 and all(argv[1:4] == ["sweep", "run", "--recipe"] for argv in calls) + assert len(backend_calls) == 2 + + +def test_downloaded_run_verifies_without_original_runtime_paths(completed, tmp_path): + original = completed[0] + relocated = tmp_path / "download" + shutil.copytree(original, relocated) + # Remove only the test-owned run tree, so no original runtime input paths exist. + shutil.rmtree(original) + assert "expected result verified" in helper.verify(relocated) + + +@pytest.mark.parametrize("mutation", ["identity", "truncated", "wrong_verdict", "backend"]) +def test_inconsistent_report_rejected(completed, mutation): + output = completed[0] + path = output / "corrected/sanitizer_report.json" + data = json.loads(path.read_text(encoding="utf-8")) + if mutation == "identity": + data["checks"][0]["kernel_results"][0]["identity"]["entry_offset"] = 0 + elif mutation == "truncated": + data["checks"][0]["diagnostics_truncated"] = True + elif mutation == "backend": + data["checks"][0]["backend"]["sha256"] = "f" * 64 + else: + data["overall_verdict"] = "warn" + data["checks"][0]["verdict"] = "warn" + data["checks"][0]["kernel_results"][0]["verdict"] = "warn" + path.write_text(json.dumps(data), encoding="utf-8") + with pytest.raises(ValueError): + helper.verify(output) + + +@pytest.mark.parametrize("foreign", ["object", "target", "image"]) +def test_foreign_diagnostic_identity_rejected_even_when_report_and_log_agree(completed, foreign): + output = completed[0] + path = output / "missing/sanitizer_report.json" + data = json.loads(path.read_text(encoding="utf-8")) + check = data["checks"][0] + identity = check["kernel_results"][0]["identity"] + original = check["findings"][0]["message"] + prefix = f"{identity['code_object']}:{identity['target']}[{identity['code_object_index']}]:" + replacements = { + "object": f"{identity['code_object']}-foreign:{identity['target']}[{identity['code_object_index']}]:", + "target": f"{identity['code_object']}:gfx999[{identity['code_object_index']}]:", + "image": f"{identity['code_object']}:{identity['target']}[{identity['code_object_index'] + 1}]:", + } + changed = original.replace(prefix, replacements[foreign], 1) + assert changed != original + check["findings"][0]["message"] = changed + check["kernel_results"][0]["findings"][0]["message"] = changed + log = output / "missing/waitcheck/waitcheck-0.log" + log.write_text(log.read_text(encoding="utf-8").replace(original, changed), encoding="utf-8") + path.write_text(json.dumps(data), encoding="utf-8") + with pytest.raises(ValueError): + helper.verify(output) + + +@pytest.mark.parametrize("label,updates", [ + ("corrected", {"reported_count": 9}), + ("missing", {"reported_count": 9, "emitted_records": 9}), + ("missing", {"requested_limit": 1}), + ("corrected", {"count_meaning": "exact", "details": "complete", "native_analysis": "complete"}), + ("missing", {"count_meaning": "exact"}), + ("mismatch", {"reported_count": 9}), +]) +def test_coherently_declared_metadata_must_agree_with_raw_control_evidence(completed, label, updates): + output = completed[0] + path = output / label / "sanitizer_report.json" + data = json.loads(path.read_text(encoding="utf-8")) + check = data["checks"][0] + entry = check["kernel_results"][0] + entry["diagnostic_summary"].update(updates) + # Keep check aggregation valid so the regression reaches the control verifier. + check["diagnostic_summary"] = aggregate_diagnostic_summaries((KernelCheckResult.from_dict(entry),)).to_dict() + path.write_text(json.dumps(data), encoding="utf-8") + with pytest.raises(ValueError, match="diagnostic"): + helper.verify(output) + + +@pytest.mark.parametrize("nullable", [False, True]) +def test_legacy_controls_allow_absent_or_null_optional_diagnostic_summary(completed, nullable): + output = completed[0] + for label in helper.LABELS: + path = output / label / "sanitizer_report.json" + data = json.loads(path.read_text(encoding="utf-8")) + check = data["checks"][0] + for item in (check, check["kernel_results"][0]): + if nullable: + item["diagnostic_summary"] = None + else: + item.pop("diagnostic_summary", None) + path.write_text(json.dumps(data), encoding="utf-8") + assert "expected result verified" in helper.verify(output) + + +def test_changed_recipe_refused_before_process(prepared, monkeypatch, tmp_path): + path = prepared / "missing.json" + data = json.loads(path.read_text(encoding="utf-8")) + data["sanitizer_plan"]["sanitizers"] = ["asan"] + path.write_text(json.dumps(data), encoding="utf-8") + monkeypatch.setattr(helper.subprocess, "Popen", lambda *args, **kwargs: pytest.fail("must not launch")) + with pytest.raises(ValueError, match="recipe changed"): + helper.run(argparse.Namespace(prepared=prepared, output=tmp_path / "run")) + + +@pytest.mark.parametrize("mutation", ["argv", "identity_path", "extra_log", "environment"]) +def test_changed_execution_scope_or_extra_output_rejected(completed, mutation): + output = completed[0] + path = output / "run.json" + manifest = json.loads(path.read_text(encoding="utf-8")) + if mutation == "argv": + manifest["commands"]["corrected"]["argv"][4] = manifest["commands"]["missing"]["argv"][4] + elif mutation == "identity_path": + manifest["identities"]["corrected"]["code_object"] = manifest["identities"]["missing"]["code_object"] + elif mutation == "environment": + manifest["environment_override"]["ROCJITSU_PREBUILT"] += "-other" + else: + log = output / "corrected/waitcheck/waitcheck-0.log" + with log.open("a", encoding="utf-8") as stream: + stream.write("unexpected analysis error\n") + path.write_text(json.dumps(manifest), encoding="utf-8") + with pytest.raises(ValueError): + helper.verify(output) + + +def test_control_cannot_be_used_as_aorta_executable(prepared, tmp_path, monkeypatch): + fake_cli = tmp_path / "control-copy" + shutil.copyfile(prepared / "inputs/missing.bin", fake_cli) + fake_cli.chmod(0o755) + bundle = tmp_path / "bundle" + (bundle / "bin").mkdir(parents=True) + scanner = bundle / "bin/rj_waitcheck" + scanner.write_bytes(b"inert scanner") + scanner.chmod(0o755) + monkeypatch.setattr(helper.subprocess, "Popen", lambda *args, **kwargs: pytest.fail("must not execute input")) + with pytest.raises(ValueError, match="control input"): + helper.run(argparse.Namespace(prepared=prepared, output=tmp_path / "run", aorta=fake_cli, + prebuilt=bundle, tool_sha256=helper.digest(scanner))) + + +def test_timeout_captures_failure_without_clean_claim(tmp_path, monkeypatch): + class Process: + killed = False + def wait(self, timeout): + if not self.killed: + raise subprocess.TimeoutExpired("fake", timeout) + return -9 + def kill(self): + self.killed = True + process = Process() + monkeypatch.setattr(helper.subprocess, "Popen", lambda *args, **kwargs: process) + monkeypatch.setattr(helper.os, "name", "nt") + row = helper.capture(["fake-aorta"], tmp_path, "timeout", {}) + assert process.killed and row["timed_out"] and row["returncode"] == -9 + assert (tmp_path / "timeout.stderr.txt").exists() + + +def test_existing_and_dangling_outputs_refused(tmp_path): + existing = tmp_path / "existing" + existing.mkdir() + (existing / "keep").write_text("unchanged", encoding="utf-8") + with pytest.raises(ValueError): + helper.fresh(existing) + assert (existing / "keep").read_text(encoding="utf-8") == "unchanged" + link, target = tmp_path / "link", tmp_path / "target" + try: + link.symlink_to(target, target_is_directory=True) + except OSError as exc: + pytest.skip(f"symlinks unavailable: {exc}") + with pytest.raises(ValueError): + helper.fresh(link) + assert not target.exists() diff --git a/tests/instrumentation/rocjitsu_sanitizers/test_waitcheck_failures.py b/tests/instrumentation/rocjitsu_sanitizers/test_waitcheck_failures.py new file mode 100644 index 000000000..fa4e85d67 --- /dev/null +++ b/tests/instrumentation/rocjitsu_sanitizers/test_waitcheck_failures.py @@ -0,0 +1,169 @@ +"""Recipe-to-report failure boundaries, using inert files and a fake executor.""" +from __future__ import annotations + +from functools import partial +import hashlib +from pathlib import Path + +import pytest +import yaml + +from aorta.instrumentation.rocjitsu_sanitizers import pipeline, waitcheck +from aorta.instrumentation.rocjitsu_sanitizers.execution import ProcessResult +from aorta.instrumentation.rocjitsu_sanitizers.recipe import execute_sanitizer_run +from aorta.instrumentation.rocjitsu_sanitizers.report import read_report +from aorta.triage.recipe import RecipeSchemaError + + +@pytest.fixture +def setup_case(tmp_path: Path, monkeypatch: pytest.MonkeyPatch): + # Prevent discovery of the developer's installed tools in every test. + monkeypatch.delenv("ROCJITSU_BUILD", raising=False) + monkeypatch.setattr(waitcheck.shutil, "which", lambda name: None) + bundle = tmp_path / "bundle" + binary = bundle / "bin" / "rj_waitcheck" + binary.parent.mkdir(parents=True) + binary.write_bytes(b"inert scanner; never execute") + monkeypatch.setenv("ROCJITSU_PREBUILT", str(bundle)) + artifact = tmp_path / "control.co" + artifact.write_bytes(b"inert code object; never execute") + data = { + "schema_version": 1, "mode": "sanitizer", "ticket": "FAILURE-BOUNDARY", + "sanitizer_plan": { + "target": "gfx1250", "source": {"kind": "kernel", "kernel": { + "name": "control", "code_object": str(artifact), + "code_object_sha256": hashlib.sha256(artifact.read_bytes()).hexdigest(), + "code_object_index": 0, "entry_offset": 256}}, + "scope": {"kind": "kernel"}, + "selection": {"requirement": "top_dispatch_count", "top_n": 1}, + "sanitizers": ["waitcheck"], + "policy": {"consan_policy": "strict", "on_missing_backend": "fail", + "timeout_seconds": 120, "waitcheck_max_diagnostics": 64}, + "output": {"report": "sanitizer_report.json"}}} + calls = [] + + def never_execute(*args, **kwargs): + pytest.fail("native executor must not be reached") + + monkeypatch.setattr(pipeline, "run_waitcheck", + partial(waitcheck.run_waitcheck, execute=never_execute)) + + def run(execute=None): + if execute is not None: + def capture(argv, *, timeout_seconds, **kwargs): + calls.append(tuple(argv)) + assert timeout_seconds == 120 + assert argv[argv.index("--max-diagnostics") + 1] == "64" + assert argv[argv.index("--kernel-entry") + 1] == "0x100" + assert argv[argv.index("--code-object-index") + 1] == "0" + return execute(argv) + monkeypatch.setattr(pipeline, "run_waitcheck", + partial(waitcheck.run_waitcheck, execute=capture)) + recipe = tmp_path / "case.yaml" + recipe.write_text(yaml.safe_dump(data), encoding="utf-8") + path = execute_sanitizer_run(recipe, output_dir=tmp_path / "out") + return read_report(path) + + return data, artifact, binary, run, calls + + +@pytest.mark.parametrize("problem,reason,kernel_state", [ + ("missing", "code_object_not_found", "not_checked"), + ("empty", "code_object_empty", "not_checked"), + ("changed", "code_object_digest_mismatch", "error"), +]) +def test_object_problem_persists_nonclean_report(setup_case, tmp_path, problem, reason, kernel_state): + _, artifact, _, run, calls = setup_case + if problem == "missing": + artifact.unlink() + else: + artifact.write_bytes(b"" if problem == "empty" else b"changed bytes") + report = run() + assert report.overall_verdict.value == "error" + assert report.execution_status.value == "error" + check = report.checks[0] + assert check.reason == "worklist_not_fully_checked" + result = check.kernel_results[0] + assert (result.state.value, result.reason, result.returncode) == (kernel_state, reason, None) + assert not calls + assert not (tmp_path / "out/waitcheck/waitcheck-0.log").exists() + + +@pytest.mark.parametrize("problem,reason", [("missing", "rj_waitcheck_not_found"), + ("empty", "rj_waitcheck_empty")]) +def test_missing_backend_fail_policy_preserves_not_checked(setup_case, problem, reason): + data, _, binary, run, calls = setup_case + assert data["sanitizer_plan"]["policy"]["on_missing_backend"] == "fail" + if problem == "missing": + binary.unlink() + else: + binary.write_bytes(b"") + report = run() + assert (report.overall_verdict.value, report.execution_status.value) == ("not_checked", "not_checked") + assert report.checks[0].reason == reason + assert report.checks[0].kernel_results[0].reason == reason + assert not calls + + +@pytest.mark.parametrize("field,value", [("entry_offset", -1), ("code_object_index", True)]) +def test_invalid_selector_rejected_before_report_or_execution(setup_case, tmp_path, field, value): + data, _, _, run, calls = setup_case + data["sanitizer_plan"]["source"]["kernel"][field] = value + with pytest.raises(RecipeSchemaError, match=field): + run() + assert not calls + assert not (tmp_path / "out/sanitizer_report.json").exists() + + +@pytest.mark.parametrize("problem,reason,state", [ + ("timeout", "waitcheck_timeout", "timed_out"), + ("invalid_entry", "waitcheck_backend_exit_2", "error"), + ("unsupported_flag", "waitcheck_backend_exit_2", "error"), + ("malformed_output", "waitcheck_diagnostics_error", "error"), +]) +def test_backend_failure_survives_recipe_report_boundary(setup_case, tmp_path, problem, reason, state): + _, _, _, run, calls = setup_case + messages = {"invalid_entry": "synthetic: requested kernel entry not found", + "unsupported_flag": "synthetic: unknown --kernel-entry option"} + + def execute(argv): + if problem == "timeout": + return ProcessResult(tuple(argv), None, "synthetic partial output", "", timed_out=True) + if problem in messages: + return ProcessResult(tuple(argv), 2, "", messages[problem]) + return ProcessResult(tuple(argv), 0, "synthetic incompatible output format", "") + + report = run(execute) + assert (report.overall_verdict.value, report.execution_status.value) == ("error", "error") + check = report.checks[0] + assert check.state.value == state + result = check.kernel_results[0] + assert result.reason.startswith(reason) + assert result.state.value == state + if problem in messages: + assert messages[problem] in result.reason + assert result.returncode == 2 + if problem == "malformed_output": + assert result.returncode == 0 # Zero native exit cannot hide an unreadable result. + assert len(calls) == 1 + raw = (tmp_path / "out/waitcheck/waitcheck-0.log").read_text(encoding="utf-8") + expected = ("synthetic partial output" if problem == "timeout" else + "synthetic incompatible output format" if problem == "malformed_output" else messages[problem]) + assert expected in raw + + +def test_capped_finding_remains_warning_with_explicit_incomplete_count(setup_case): + _, artifact, _, run, calls = setup_case + prefix = f"{artifact}:gfx1250[0]" + output = (f"rj_waitcheck: {prefix}:kernel=.text+0x100: " + "instructions=4 memory-events=2 diagnostics=>=1\n" + f"{prefix}:.text+0x108: missing s_wait_kmcnt <= 0 before use of s4\n") + report = run(lambda argv: ProcessResult(tuple(argv), 4, output, "")) + assert (report.overall_verdict.value, report.execution_status.value) == ("warn", "complete") + check = report.checks[0] + assert check.state.value == "ran" + assert check.diagnostics_truncated is True + assert check.kernel_results[0].diagnostics_truncated is True + assert check.kernel_results[0].returncode == 4 + assert len(check.findings) == 1 + assert len(calls) == 1 diff --git a/tests/instrumentation/rocjitsu_sanitizers/test_waitcheck_isolation_recipe.py b/tests/instrumentation/rocjitsu_sanitizers/test_waitcheck_isolation_recipe.py new file mode 100644 index 000000000..7e6349df8 --- /dev/null +++ b/tests/instrumentation/rocjitsu_sanitizers/test_waitcheck_isolation_recipe.py @@ -0,0 +1,78 @@ +"""Check draft two-kernel recipes without compiling, launching or scanning code.""" +from __future__ import annotations + +import hashlib +import os +from pathlib import Path +import subprocess + +import pytest +import yaml + +from aorta.instrumentation.rocjitsu_sanitizers.models import KernelIdentity +from aorta.instrumentation.rocjitsu_sanitizers.recipe import load_sanitizer_recipe +from aorta.instrumentation.rocjitsu_sanitizers.waitcheck import waitcheck_argv +from aorta.triage.recipe import RecipeSchemaError + +ASSETS = Path(__file__).resolve().parents[3] / "examples/sanitizers/waitcheck/assets" + + +@pytest.fixture(autouse=True) +def prohibit_execution(monkeypatch): + def forbidden(*args, **kwargs): + raise AssertionError("these recipe checks must not launch any process") + monkeypatch.setattr(subprocess, "Popen", forbidden) + monkeypatch.setattr(os, "system", forbidden) + + +@pytest.mark.parametrize("label", ["missing", "corrected"]) +@pytest.mark.parametrize("resolved_index", [False, True]) +def test_unresolved_isolation_selector_is_refused(tmp_path, label, resolved_index): + data = yaml.safe_load((ASSETS / f"waitcheck-isolation-{label}.yaml").read_text(encoding="utf-8")) + if resolved_index: + data["sanitizer_plan"]["source"]["kernel"]["code_object_index"] = 0 + path = tmp_path / "unresolved.yaml" + path.write_text(yaml.safe_dump(data), encoding="utf-8") + with pytest.raises(RecipeSchemaError, match="entry_offset" if resolved_index else "code_object_index"): + load_sanitizer_recipe(path) + + +def test_prepared_pair_selects_distinct_entries_in_same_object(tmp_path): + # Inert local data and synthetic selectors test the recipe/argv contract only. + # These offsets are NOT proposed entry values for the HIP source. + artifact = tmp_path / "two-kernel-fixture.co" + artifact.write_bytes(b"inert shared-object recipe identity fixture") + fingerprint = hashlib.sha256(artifact.read_bytes()).hexdigest() + identities, argv_by_label = {}, {} + for label, entry in (("missing", 0), ("corrected", 0x80)): + data = yaml.safe_load((ASSETS / f"waitcheck-isolation-{label}.yaml").read_text(encoding="utf-8")) + spec = data["sanitizer_plan"]["source"]["kernel"] + spec.update(code_object=artifact.name, code_object_sha256=fingerprint, + code_object_index=2, entry_offset=entry) + path = tmp_path / f"{label}.yaml" + path.write_text(yaml.safe_dump(data), encoding="utf-8") + loaded = load_sanitizer_recipe(path) + assert loaded.target == "gfx1250" + assert loaded.sanitizers == ("waitcheck",) + assert loaded.scope_kind == "kernel" and loaded.top_n == 1 + assert loaded.timeout_seconds == 120 and loaded.waitcheck_max_diagnostics == 64 + assert len(loaded.kernel_specs) == 1 + kernel = loaded.kernel_specs[0] + identity = KernelIdentity(target=loaded.target, **kernel.to_entry()) + assert identity.name == f"waitcheck_{label}" + assert identity.exact and not identity.code_object_scan + assert identity.entry_offset == entry + identities[label] = identity + command = waitcheck_argv(Path("rj_waitcheck"), identity, max_diagnostics=loaded.waitcheck_max_diagnostics) + assert command[1] == str(artifact.resolve()) + assert command[command.index("--code-object-index") + 1] == "2" + assert command[command.index("--kernel-entry") + 1] == hex(entry) + assert command[command.index("--max-diagnostics") + 1] == "64" + argv_by_label[label] = command + missing, corrected = identities["missing"], identities["corrected"] + assert missing.code_object == corrected.code_object + assert missing.code_object_sha256 == corrected.code_object_sha256 == fingerprint + assert missing.code_object_index == corrected.code_object_index == 2 + assert missing.entry_offset != corrected.entry_offset + assert missing.stable_key != corrected.stable_key + assert argv_by_label["missing"] != argv_by_label["corrected"] diff --git a/tests/instrumentation/rocjitsu_sanitizers/test_waitcheck_native_cases.py b/tests/instrumentation/rocjitsu_sanitizers/test_waitcheck_native_cases.py new file mode 100644 index 000000000..6c3f73bfe --- /dev/null +++ b/tests/instrumentation/rocjitsu_sanitizers/test_waitcheck_native_cases.py @@ -0,0 +1,78 @@ +"""Local schema/argv checks for draft selector and cap cases; no native execution.""" +from __future__ import annotations + +from copy import deepcopy +import hashlib +import os +from pathlib import Path +import subprocess + +import pytest +import yaml + +from aorta.instrumentation.rocjitsu_sanitizers.models import KernelIdentity +from aorta.instrumentation.rocjitsu_sanitizers.recipe import load_sanitizer_recipe +from aorta.instrumentation.rocjitsu_sanitizers.waitcheck import waitcheck_argv +from aorta.triage.recipe import RecipeSchemaError + +ASSETS = Path(__file__).resolve().parents[3] / "examples/sanitizers/waitcheck/assets" +CASES = ("invalid-image", "invalid-entry", "cap-high", "cap-low") + + +@pytest.fixture(autouse=True) +def no_execution(monkeypatch): + def forbidden(*args, **kwargs): + raise AssertionError("recipe checks cannot launch a process") + monkeypatch.setattr(subprocess, "Popen", forbidden) + monkeypatch.setattr(os, "system", forbidden) + + +@pytest.mark.parametrize("case", CASES) +@pytest.mark.parametrize("resolved_image", [False, True]) +def test_unresolved_identity_is_not_bypassed(tmp_path, case, resolved_image): + data = yaml.safe_load((ASSETS / f"waitcheck-{case}.yaml").read_text(encoding="utf-8")) + if resolved_image: + data["sanitizer_plan"]["source"]["kernel"]["code_object_index"] = 0 + path = tmp_path / "draft.yaml" + path.write_text(yaml.safe_dump(data), encoding="utf-8") + with pytest.raises(RecipeSchemaError, match="entry_offset" if resolved_image else "code_object_index"): + load_sanitizer_recipe(path) + + +@pytest.mark.parametrize("case,image,entry,cap", [ + ("invalid-image", 1, 256, 64), + ("invalid-entry", 0, 20480, 64), + ("cap-high", 0, 128, 64), + ("cap-low", 0, 128, 1), +]) +def test_resolved_recipe_preserves_explicit_scan_scope(tmp_path, case, image, entry, cap): + # Selectors here test schema/argv only, not whether a kernel really exists. + artifact = tmp_path / "inert.co" + artifact.write_bytes(b"inert local recipe fixture") + data = yaml.safe_load((ASSETS / f"waitcheck-{case}.yaml").read_text(encoding="utf-8")) + data["sanitizer_plan"]["source"]["kernel"].update( + code_object=artifact.name, code_object_sha256=hashlib.sha256(artifact.read_bytes()).hexdigest(), + code_object_index=image, entry_offset=entry, + ) + path = tmp_path / "prepared.yaml" + path.write_text(yaml.safe_dump(data), encoding="utf-8") + recipe = load_sanitizer_recipe(path) + assert recipe.target == "gfx1250" and recipe.sanitizers == ("waitcheck",) + assert recipe.scope_kind == "kernel" and recipe.top_n == 1 + assert recipe.timeout_seconds == 120 and recipe.waitcheck_max_diagnostics == cap + assert len(recipe.kernel_specs) == 1 + identity = KernelIdentity(target=recipe.target, **recipe.kernel_specs[0].to_entry()) + assert identity.exact and not identity.code_object_scan + argv = waitcheck_argv(Path("rj_waitcheck"), identity, max_diagnostics=cap) + assert argv[argv.index("--code-object-index") + 1] == str(image) + assert argv[argv.index("--kernel-entry") + 1] == hex(entry) + assert argv[argv.index("--max-diagnostics") + 1] == str(cap) + + +def test_cap_comparison_changes_only_cap_and_metadata(): + high = yaml.safe_load((ASSETS / "waitcheck-cap-high.yaml").read_text(encoding="utf-8")) + low = yaml.safe_load((ASSETS / "waitcheck-cap-low.yaml").read_text(encoding="utf-8")) + normalized = deepcopy(low) + normalized["ticket"], normalized["description"] = high["ticket"], high["description"] + normalized["sanitizer_plan"]["policy"]["waitcheck_max_diagnostics"] = 64 + assert normalized == high diff --git a/tests/instrumentation/rocjitsu_sanitizers/test_waitcheck_report.py b/tests/instrumentation/rocjitsu_sanitizers/test_waitcheck_report.py new file mode 100644 index 000000000..977564f1b --- /dev/null +++ b/tests/instrumentation/rocjitsu_sanitizers/test_waitcheck_report.py @@ -0,0 +1,661 @@ +"""Offline reader tests: arbitrary synthetic entries; no native execution.""" +import importlib.util +import json +from pathlib import Path + +import pytest + +from aorta.instrumentation.rocjitsu_sanitizers.models import ( + CheckResult, ExecutionState, KernelCheckResult, KernelIdentity, KernelObservation, + KernelWorklist, SelectionRequirement, Verdict, + aggregate_diagnostic_summaries) +from aorta.instrumentation.rocjitsu_sanitizers.report import build_report, write_report +from aorta.instrumentation.rocjitsu_sanitizers.waitcheck import parse_waitcheck_text + +ROOT = Path(__file__).resolve().parents[3] +SPEC = importlib.util.spec_from_file_location( + "waitcheck_report", ROOT / '.agents/skills/rocm-sanitizer-triage/scripts/waitcheck_report.py') +helper = importlib.util.module_from_spec(SPEC) +SPEC.loader.exec_module(helper) + + +@pytest.fixture +def evidence(tmp_path): + identity = KernelIdentity(name='user_attention_variant_42', target='gfx950', + code_object='/user/project/objects/attention.so', + code_object_sha256='a' * 64, code_object_index=7, entry_offset=0x240) + worklist = KernelWorklist(requirement=SelectionRequirement.TOP_DISPATCH_COUNT, top_n=1, + kernels=(KernelObservation(identity=identity, dispatch_count=1, sources=('synthetic',)),)) + + def create(kind='clean'): + prefix = f'{identity.code_object}:gfx950[7]' + hazard = kind in ('warn', 'cap') + log = (f'rj_waitcheck: {prefix}:kernel=.text+0x240: instructions=4 memory-events=2 ' + f'diagnostics={">=" if kind == "cap" else ""}{1 if hazard else 0}\n') + if hazard: + log += f'{prefix}:.text+0x248: missing s_waitcnt lgkmcnt(0)\n' + parsed = parse_waitcheck_text(log, expected=identity) + verdict = Verdict.WARN if hazard else Verdict.PASS + result = KernelCheckResult(identity=identity, state=ExecutionState.RAN, verdict=verdict, + returncode=4 if hazard else 0, findings=parsed.findings, + diagnostics_truncated=kind == 'cap') + check = CheckResult(sanitizer='waitcheck', state=ExecutionState.RAN, verdict=verdict, + findings=result.findings, kernel_results=(result,), + diagnostics_truncated=kind == 'cap', backend=(('path', '/tools/rj_waitcheck'),)) + if kind == 'digest': + result = KernelCheckResult(identity=identity, state=ExecutionState.ERROR, + verdict=Verdict.ERROR, reason='code_object_digest_mismatch') + check = CheckResult(sanitizer='waitcheck', state=ExecutionState.ERROR, + verdict=Verdict.ERROR, reason='worklist_not_fully_checked', kernel_results=(result,)) + elif kind == 'timeout': + result = KernelCheckResult(identity=identity, state=ExecutionState.TIMED_OUT, + verdict=Verdict.ERROR, reason='waitcheck_timeout') + check = CheckResult(sanitizer='waitcheck', state=ExecutionState.TIMED_OUT, + verdict=Verdict.ERROR, reason='worklist_not_fully_checked', kernel_results=(result,)) + log = 'partial stdout before timeout\n' + elif kind == 'dry': + check = CheckResult(sanitizer='waitcheck', state=ExecutionState.NOT_CHECKED, + verdict=Verdict.NOT_CHECKED, reason='dry_run') + report_path, log_path = tmp_path / 'report.json', tmp_path / 'native.log' + write_report(build_report(target=identity.target, worklist=worklist, checks=(check,)), report_path) + log_path.write_text(log, encoding='utf-8') + return report_path, log_path + return create + + +@pytest.mark.parametrize('kind', ['clean', 'warn', 'cap']) +def test_arbitrary_identity_and_reported_result_with_limits(evidence, kind): + report, log = evidence(kind) + summary, code = helper.summarize(report, log) + assert code == 0 + assert 'Raw summary agrees' in summary + assert 'user_attention_variant_42' in summary + assert 'not independently verified provenance' in summary + assert 'AORTA CLI exit is not recorded' in summary + if kind in ('warn', 'cap'): + assert '"native_returncode": 4' in summary + assert '"overall_verdict": "warn"' in summary + if kind == 'cap': + assert 'lower bound' in summary + + +def test_missing_raw_is_a_limit_not_evidence_of_clean(evidence): + report, _ = evidence() + summary, code = helper.summarize(report) + assert code == 0 + assert 'No raw log supplied' in summary + assert 'Raw findings/entry evidence was not compared' in summary + + +@pytest.mark.parametrize('kind,reason', [('digest', 'code_object_digest_mismatch'), ('dry', 'dry_run')]) +def test_prelaunch_outcome_needs_no_raw(evidence, kind, reason): + report, _ = evidence(kind) + summary, code = helper.summarize(report) + assert code == 0 + assert reason in summary + assert 'No completed single-entry check is established' in summary + + +def test_timeout_keeps_partial_log_without_parse_pass(evidence): + report, log = evidence('timeout') + summary, code = helper.summarize(report, log) + assert code == 0 + assert 'waitcheck_timeout' in summary + assert 'partial stdout before timeout' in summary + assert 'not parsed as a completed scan' in summary + + +@pytest.mark.parametrize('mutation', ['path', 'entry', 'count', 'extra', 'unsupported', 'truncation']) +def test_raw_contradictions_and_unknown_lines_do_not_confirm_clean(evidence, mutation): + report, log = evidence() + raw = log.read_text() + if mutation == 'path': + raw = raw.replace('/user/project/', '/another/') + elif mutation == 'entry': + raw = raw.replace('0x240', '0x280') + elif mutation == 'count': + raw = raw.replace('diagnostics=0', 'diagnostics=2') + elif mutation == 'extra': + raw += 'ERROR: incomplete instruction support\n' + elif mutation == 'truncation': + raw = raw.replace('diagnostics=0', 'diagnostics=>=0') + else: + raw = 'new incompatible summary vocabulary\n' + log.write_text(raw) + summary, code = helper.summarize(report, log) + assert code == 2 + assert 'INCONSISTENT OR UNSUPPORTED' in summary + assert 'Raw summary agrees' not in summary + + +@pytest.mark.parametrize('mutation', ['scope', 'identity', 'unknown', 'invalid']) +def test_invalid_or_unsupported_report_refused(evidence, mutation): + report, _ = evidence() + data = json.loads(report.read_text()) + if mutation == 'scope': + data['checks'].append(data['checks'][0]) + elif mutation == 'identity': + data['checks'][0]['kernel_results'][0]['identity']['entry_offset'] = 0 + elif mutation == 'unknown': + data['checks'][0]['completion_status'] = 'unsupported-instructions' + else: + data['overall_verdict'] = 'error' + report.write_text(json.dumps(data)) + with pytest.raises((ValueError, TypeError)): + helper.summarize(report) + + +def test_cli_output_exclusivity_and_reader_exit_two(evidence, tmp_path): + report, log = evidence() + output = tmp_path / 'summary.md' + output.write_text('keep') + argv = ['--report', str(report), '--log', str(log), '--output', str(output)] + assert helper.main(argv) == 1 + assert output.read_text() == 'keep' + output.unlink() + log.write_text('unsupported output') + assert helper.main(argv) == 2 + assert 'INCONSISTENT OR UNSUPPORTED' in output.read_text() + + +def test_bounded_report_input(evidence, monkeypatch): + report, _ = evidence() + monkeypatch.setattr(helper, 'LIMIT', 10) + with pytest.raises(ValueError, match='exceeds'): + helper.summarize(report) + + +def test_duplicate_json_field_refused(evidence): + report, _ = evidence() + raw = report.read_text() + raw = raw.replace('"overall_verdict": "pass"', '"overall_verdict": "error", "overall_verdict": "pass"') + report.write_text(raw) + with pytest.raises(ValueError, match='duplicate JSON field'): + helper.summarize(report) + + +def test_foreign_diagnostic_is_not_attributed_to_selected_entry(evidence): + report, log = evidence('warn') + original = log.read_text().splitlines() + original[1] = original[1].replace('/user/project/', '/foreign/') + log.write_text('\n'.join(original)) + data = json.loads(report.read_text()) + for findings in (data['checks'][0]['findings'], data['checks'][0]['kernel_results'][0]['findings']): + findings[0]['message'] = original[1] + report.write_text(json.dumps(data)) + summary, code = helper.summarize(report, log) + assert code == 2 + assert 'diagnostic object/target/image prefix' in summary + + +@pytest.mark.parametrize('mutation', ['verdict', 'returncode']) +def test_aggregate_result_conflict_preserved_and_flagged(evidence, mutation): + report, log = evidence() + data = json.loads(report.read_text()) + if mutation == 'verdict': + data['overall_verdict'] = data['checks'][0]['verdict'] = 'warn' + else: + data['checks'][0]['returncode'] = 4 + report.write_text(json.dumps(data)) + summary, code = helper.summarize(report, log) + assert code == 2 + assert 'Report consistency: INCONSISTENT' in summary + assert '"native_returncode": 0' in summary + + +@pytest.mark.parametrize('kind,native_exit', [('clean', 4), ('warn', 0), ('clean', None), ('warn', None)]) +def test_report_only_native_completion_conflict(evidence, kind, native_exit): + report, _ = evidence(kind) + data = json.loads(report.read_text()) + data['checks'][0]['kernel_results'][0]['returncode'] = native_exit + report.write_text(json.dumps(data)) + summary, code = helper.summarize(report) + assert code == 2 + assert 'Report consistency: INCONSISTENT' in summary + assert 'native exit is missing or contradicts' in summary + assert f'"native_returncode": {json.dumps(native_exit)}' in summary + assert f'"entry_verdict": "{"pass" if kind == "clean" else "warn"}"' in summary + + +def rewrite_evidence(report_path, log_path, text, *, include_metadata=False, requested_limit=None): + """Portable synthetic derivatives; never execute a native program.""" + from dataclasses import replace + from aorta.instrumentation.rocjitsu_sanitizers.report import read_report + report = read_report(report_path) + previous = report.checks[0].kernel_results[0] + parsed = parse_waitcheck_text(text, expected=previous.identity) + metadata = replace(parsed.diagnostic_summary, requested_limit=requested_limit) if include_metadata else None + result = replace(previous, findings=parsed.findings, diagnostics_truncated=parsed.diagnostics_truncated, + diagnostic_summary=metadata) + check = replace(report.checks[0], findings=result.findings, kernel_results=(result,), + diagnostics_truncated=result.diagnostics_truncated, + diagnostic_summary=aggregate_diagnostic_summaries((result,)) if include_metadata else None) + write_report(build_report(target=report.target, worklist=report.worklist, checks=(check,)), report_path) + log_path.write_text(text, encoding='utf-8') + + +def test_duplicate_records_are_not_omitted_details(evidence): + report, log = evidence('warn') + text = log.read_text().replace('diagnostics=1', 'diagnostics=2') + text += text.splitlines()[1] + '\n' + rewrite_evidence(report, log, text, include_metadata=True) + summary, code = helper.summarize(report, log) + assert code == 0 + assert 'emitted records=2, retained unique findings=1' in summary + assert 'compatibility gap' not in summary + assert 'native analysis=unknown' in summary + + +@pytest.mark.parametrize('include_metadata', [False, True]) +def test_plain_limited_details_remain_explicit_compatibility_gap(evidence, include_metadata): + report, log = evidence('warn') + text = log.read_text().replace('diagnostics=1', 'diagnostics=3') + rewrite_evidence(report, log, text, include_metadata=include_metadata, requested_limit=1) + summary, code = helper.summarize(report, log) + assert code == 2 + assert 'compatibility gap' in summary + assert 'reported diagnostics=3, emitted records=1, retained unique findings=1' in summary + assert 'count meaning=unknown' in summary + assert 'native analysis=unknown' in summary + assert '"native_returncode": 4' in summary + assert '"diagnostics_truncated": false' in summary + assert '"reported_count": 3' in summary + assert '"emitted_records": 1' in summary + + +@pytest.mark.parametrize('include_metadata', [False, True]) +def test_explicit_storage_omission_is_supported_separately_from_completion(evidence, include_metadata): + report, log = evidence('warn') + text = log.read_text().replace('diagnostics=1', 'diagnostics=3') + text += 'omitted 2 diagnostic(s) after limit\n' + rewrite_evidence(report, log, text, include_metadata=include_metadata, requested_limit=1) + summary, code = helper.summarize(report, log) + assert code == 0 + assert 'Count meaning=exact; details=incomplete; native analysis=complete' in summary + assert '"diagnostics_truncated": false' in summary + + +@pytest.mark.parametrize('signal', ['omitted at least 2 diagnostic(s) after limit', + 'analysis stopped early; additional diagnostics may exist']) +def test_explicit_early_stop_is_not_completed_analysis(evidence, signal): + report, log = evidence('cap') + text = log.read_text().replace('instructions=4', 'instructions=>=4').replace('memory-events=2', 'memory-events=>=2') + if signal.startswith('omitted'): + text = text.replace('diagnostics=>=1', 'diagnostics=>=3') + text += signal + '\n' + rewrite_evidence(report, log, text, include_metadata=True, requested_limit=1) + summary, code = helper.summarize(report, log) + assert code == 0 + assert 'Count meaning=lower_bound' in summary + assert 'native analysis=incomplete' in summary + assert '"execution_status": "complete"' in summary + + +def test_reaching_requested_cap_does_not_establish_completeness(evidence): + report, log = evidence('warn') + rewrite_evidence(report, log, log.read_text(), include_metadata=True, requested_limit=1) + summary, code = helper.summarize(report, log) + assert code == 0 + assert '"requested_limit": 1' in summary + assert 'Count meaning=unknown; details=unknown; native analysis=unknown' in summary + + +def test_legacy_report_without_new_or_old_optional_metadata_is_readable(evidence): + report, log = evidence('clean') + data = json.loads(report.read_text()) + for item in (data['checks'][0], data['checks'][0]['kernel_results'][0]): + item.pop('diagnostics_truncated') + item.pop('diagnostic_summary', None) + report.write_text(json.dumps(data)) + summary, code = helper.summarize(report, log) + assert code == 0 + assert 'Legacy report' in summary + assert 'Count meaning=unknown; details=unknown; native analysis=unknown' in summary + + +def test_report_only_keeps_metadata_and_missing_raw_limit(evidence): + report, log = evidence('warn') + rewrite_evidence(report, log, log.read_text(), include_metadata=True, requested_limit=64) + summary, code = helper.summarize(report) + assert code == 0 + assert '"requested_limit": 64' in summary + assert '"native_analysis": "unknown"' in summary + assert 'Raw findings/entry evidence was not compared' in summary + assert '"diagnostic_summary_raw": null' in summary + + +@pytest.mark.parametrize('key,value', [('reported_count', True), ('emitted_records', -1), + ('requested_limit', 0), ('details', 'yes'), + ('native_analysis', 'success'), ('unknown_key', 1)]) +def test_malformed_new_metadata_is_rejected(evidence, key, value): + report, log = evidence('warn') + rewrite_evidence(report, log, log.read_text(), include_metadata=True, requested_limit=1) + data = json.loads(report.read_text()) + data['checks'][0]['kernel_results'][0]['diagnostic_summary'][key] = value + report.write_text(json.dumps(data)) + with pytest.raises((ValueError, TypeError)): + helper.summarize(report, log) + + +def test_new_metadata_cannot_claim_a_different_raw_count(evidence): + report, log = evidence('warn') + rewrite_evidence(report, log, log.read_text(), include_metadata=True, requested_limit=1) + data = json.loads(report.read_text()) + data['checks'][0]['kernel_results'][0]['diagnostic_summary']['reported_count'] = 2 + report.write_text(json.dumps(data)) + summary, code = helper.summarize(report, log) + assert code == 2 + assert 'raw reported_count differs from report metadata' in summary + + +def test_report_completion_claim_requires_supporting_raw_signal(evidence): + report, log = evidence('warn') + rewrite_evidence(report, log, log.read_text(), include_metadata=True, requested_limit=1) + data = json.loads(report.read_text()) + for item in (data['checks'][0], data['checks'][0]['kernel_results'][0]): + item['diagnostic_summary'].update(reported_count=1, count_meaning='exact', native_analysis='complete') + report.write_text(json.dumps(data)) + summary, code = helper.summarize(report, log) + assert code == 2 + assert 'raw count_meaning does not support report metadata' in summary + + +@pytest.mark.parametrize('terminal', ['omitted 9 diagnostic(s) after limit', + 'omitted some diagnostics after limit', + 'scan complete']) +def test_conflicting_or_unknown_terminal_lines_do_not_confirm_completion(evidence, terminal): + report, log = evidence('warn') + log.write_text(log.read_text() + terminal + '\n') + summary, code = helper.summarize(report, log) + assert code == 2 + assert 'INCONSISTENT OR UNSUPPORTED' in summary + + +def test_empty_optional_metadata_stays_unknown(evidence): + report, _ = evidence('clean') + data = json.loads(report.read_text()) + for item in (data['checks'][0], data['checks'][0]['kernel_results'][0]): + item['diagnostic_summary'] = {} + report.write_text(json.dumps(data)) + summary, code = helper.summarize(report) + assert code == 0 + assert '"reported_count": null' in summary + assert '"native_analysis": "unknown"' in summary + + +@pytest.mark.parametrize('kind,count,records', [ + ('clean', 3, 0), ('clean', None, 1), ('warn', 0, None), ('warn', 1, 0), + ('warn', 1, 2), +]) +def test_report_only_metadata_count_records_and_native_exit_conflicts(evidence, kind, count, records): + report, _ = evidence(kind) + data = json.loads(report.read_text()) + data['checks'][0]['kernel_results'][0]['diagnostic_summary'] = { + 'reported_count': count, 'emitted_records': records, + } + if kind == 'warn' and records == 0: + data['checks'][0]['findings'] = [] + data['checks'][0]['kernel_results'][0]['findings'] = [] + report.write_text(json.dumps(data)) + # The reader preserves original claims and flags them without needing raw. + summary, code = helper.summarize(report) + assert code == 2 + assert 'Report consistency: INCONSISTENT' in summary + assert '"native_returncode": ' + ('0' if kind == 'clean' else '4') in summary + + +@pytest.mark.parametrize('cap,records,meaning,native', [ + (64, 3, 'unknown', 'unknown'), (1, 1, 'exact', 'complete')]) +@pytest.mark.parametrize('include_metadata', [False, True]) +def test_saved_text_derivatives_keep_the_real_omission_signal(tmp_path, cap, records, meaning, native, include_metadata): + from dataclasses import replace + text = (Path(__file__).parent / 'fixtures' / f'waitcheck-cap{cap}.txt').read_text() + identity = KernelIdentity(name='waitcheck_multiple_missing', target='gfx1250', + code_object='/fixtures/multiple_missing_waits', + code_object_sha256='a' * 64, code_object_index=0, entry_offset=0) + worklist = KernelWorklist(requirement=SelectionRequirement.TOP_TIME, top_n=1, + kernels=(KernelObservation(identity=identity, sources=('fixture',)),)) + parsed = parse_waitcheck_text(text, expected=identity) + result = KernelCheckResult(identity=identity, state=ExecutionState.RAN, verdict=Verdict.WARN, + findings=parsed.findings, returncode=4, + diagnostic_summary=replace(parsed.diagnostic_summary, requested_limit=cap) if include_metadata else None) + check = CheckResult(sanitizer='waitcheck', state=ExecutionState.RAN, verdict=Verdict.WARN, + findings=result.findings, kernel_results=(result,), + diagnostic_summary=aggregate_diagnostic_summaries((result,)) if include_metadata else None) + report_path, log_path = tmp_path / 'report.json', tmp_path / 'native.txt' + write_report(build_report(target='gfx1250', worklist=worklist, checks=(check,)), report_path) + log_path.write_text(text) + summary, code = helper.summarize(report_path, log_path) + assert code == 0 + assert f'emitted records={records}, retained unique findings={records}' in summary + assert f'Count meaning={meaning}' in summary + assert f'native analysis={native}' in summary + assert '"native_returncode": 4' in summary + if cap == 1: + assert 'Supported raw signal: omitted 2 diagnostic(s) after limit' in summary + + +@pytest.mark.parametrize('kind', ['warn', 'cap']) +@pytest.mark.parametrize('metadata', ['absent', 'empty']) +def test_report_only_hazard_exit_requires_retained_findings(evidence, kind, metadata): + report, _ = evidence(kind) + data = json.loads(report.read_text()) + for item in (data['checks'][0], data['checks'][0]['kernel_results'][0]): + item['findings'] = [] + if metadata == 'empty': + item['diagnostic_summary'] = {} + else: + item.pop('diagnostic_summary', None) + report.write_text(json.dumps(data)) + summary, code = helper.summarize(report) + assert code == 2 + assert 'native hazard exit has no retained findings' in summary + assert '"native_returncode": 4' in summary + assert '"entry_verdict": "warn"' in summary + + +@pytest.mark.parametrize('check_present,entry_present', [(False, False), (True, False), (False, True), (True, True)]) +def test_legacy_truncation_presence_is_not_a_false_claim(evidence, check_present, entry_present): + report, log = evidence('cap') + data = json.loads(report.read_text()) + for item, present in ((data['checks'][0], check_present), + (data['checks'][0]['kernel_results'][0], entry_present)): + if not present: + item.pop('diagnostics_truncated') + report.write_text(json.dumps(data)) + summary, code = helper.summarize(report, log) + assert code == 0 + assert 'Raw summary agrees' in summary + assert 'diagnostics=at least 1' in summary + assert '"diagnostics_truncated": ' + ('true' if check_present else 'null') in summary + assert '"entry_diagnostics_truncated": ' + ('true' if entry_present else 'null') in summary + + +@pytest.mark.parametrize('location', ['check', 'entry']) +def test_explicit_false_truncation_claim_still_conflicts_with_raw(evidence, location): + report, log = evidence('cap') + data = json.loads(report.read_text()) + check, entry = data['checks'][0], data['checks'][0]['kernel_results'][0] + check.pop('diagnostics_truncated') + entry.pop('diagnostics_truncated') + (check if location == 'check' else entry)['diagnostics_truncated'] = False + report.write_text(json.dumps(data)) + summary, code = helper.summarize(report, log) + assert code == 2 + assert 'truncation differs' in summary + + +@pytest.mark.parametrize('location', ['check', 'entry']) +def test_explicit_true_truncation_claim_still_conflicts_with_raw(evidence, location): + report, log = evidence('clean') + data = json.loads(report.read_text()) + check, entry = data['checks'][0], data['checks'][0]['kernel_results'][0] + check.pop('diagnostics_truncated') + entry.pop('diagnostics_truncated') + (check if location == 'check' else entry)['diagnostics_truncated'] = True + report.write_text(json.dumps(data)) + summary, code = helper.summarize(report, log) + assert code == 2 + assert 'truncation differs' in summary + + +@pytest.mark.parametrize('check_null,entry_null', [(True, False), (False, True), (True, True)]) +@pytest.mark.parametrize('with_log', [False, True]) +def test_supported_nullable_diagnostic_summary_is_missing_evidence(evidence, check_null, entry_null, with_log): + report, log = evidence('warn') + data = json.loads(report.read_text()) + for item, include in ((data['checks'][0], check_null), + (data['checks'][0]['kernel_results'][0], entry_null)): + if include: + item['diagnostic_summary'] = None + report.write_text(json.dumps(data)) + summary, code = helper.summarize(report, log if with_log else None) + assert code == 0 + assert '"entry_diagnostic_summary_reported": null' in summary + assert '"check_diagnostic_summary_reported": null' in summary + assert '"native_returncode": 4' in summary + + +@pytest.mark.parametrize('location', ['root', 'check', 'entry', 'identity']) +def test_unknown_null_fields_are_not_accepted_as_nullable_metadata(evidence, location): + report, _ = evidence() + data = json.loads(report.read_text()) + target = {'root': data, 'check': data['checks'][0], + 'entry': data['checks'][0]['kernel_results'][0], + 'identity': data['checks'][0]['kernel_results'][0]['identity']}[location] + target['unsupported_null_field'] = None + report.write_text(json.dumps(data)) + with pytest.raises(ValueError, match='unsupported fields'): + helper.summarize(report) + + +@pytest.mark.parametrize('location', ['root', 'identity']) +def test_nullable_summary_is_not_allowed_at_unmodeled_locations(evidence, location): + report, _ = evidence() + data = json.loads(report.read_text()) + target = data if location == 'root' else data['checks'][0]['kernel_results'][0]['identity'] + target['diagnostic_summary'] = None + report.write_text(json.dumps(data)) + with pytest.raises(ValueError, match='unsupported fields'): + helper.summarize(report) + + +def test_early_stop_summary_preserves_all_lower_bound_counters(evidence): + report, log = evidence('cap') + text = log.read_text().replace('instructions=4', 'instructions=>=4').replace('memory-events=2', 'memory-events=>=2') + text += 'analysis stopped early; additional diagnostics may exist\n' + rewrite_evidence(report, log, text, include_metadata=True, requested_limit=1) + summary, code = helper.summarize(report, log) + assert code == 0 + assert 'instructions=at least 4, memory-events=at least 2, diagnostics=at least 1' in summary + assert 'native analysis=incomplete' in summary + + +@pytest.mark.parametrize('nullable', [False, True]) +@pytest.mark.parametrize('early_stop', [False, True]) +def test_legacy_clean_exit_rejects_positive_raw_count_with_all_details_omitted(evidence, nullable, early_stop): + report, log = evidence('clean') + data = json.loads(report.read_text()) + for item in (data['checks'][0], data['checks'][0]['kernel_results'][0]): + item.pop('diagnostics_truncated') + if nullable: + item['diagnostic_summary'] = None + report.write_text(json.dumps(data)) + raw = log.read_text().replace('diagnostics=0', 'diagnostics=3') + if early_stop: + raw = raw.replace('instructions=4', 'instructions=>=4').replace('memory-events=2', 'memory-events=>=2') + raw = raw.replace('diagnostics=3', 'diagnostics=>=3') + raw += f'omitted {"at least " if early_stop else ""}3 diagnostic(s) after limit\n' + log.write_text(raw) + summary, code = helper.summarize(report, log) + assert code == 2 + assert 'INCONSISTENT OR UNSUPPORTED RAW EVIDENCE' in summary + assert 'native clean exit contradicts positive raw diagnostic count' in summary + assert 'Raw summary agrees' not in summary + assert '"entry_verdict": "pass"' in summary # Preserve the supplied status. + + +@pytest.mark.parametrize('nullable', [False, True]) +def test_report_only_native_zero_rejects_retained_findings_without_metadata(evidence, nullable): + report, _ = evidence('warn') + data = json.loads(report.read_text()) + check = data['checks'][0] + entry = check['kernel_results'][0] + # PASS/findings is already rejected by the model; FAIL is schema-valid and + # must still honor Waitcheck's native exit/diagnostic contract. + data['overall_verdict'] = check['verdict'] = entry['verdict'] = 'fail' + entry['returncode'] = 0 + if nullable: + check['diagnostic_summary'] = entry['diagnostic_summary'] = None + report.write_text(json.dumps(data)) + summary, code = helper.summarize(report) + assert code == 2 + assert 'native clean exit has retained findings' in summary + assert '"entry_verdict": "fail"' in summary + + +@pytest.mark.parametrize('native_exit', [2, -9]) +def test_completed_fail_rejects_unsupported_native_exit_without_raw(evidence, native_exit): + report, _ = evidence('warn') + data = json.loads(report.read_text()) + check = data['checks'][0] + entry = check['kernel_results'][0] + data['overall_verdict'] = check['verdict'] = entry['verdict'] = 'fail' + check['returncode'] = entry['returncode'] = native_exit + report.write_text(json.dumps(data)) + summary, code = helper.summarize(report) + assert code == 2 + assert 'native exit is missing or contradicts' in summary + assert '"entry_verdict": "fail"' in summary + assert f'"native_returncode": {native_exit}' in summary + + +@pytest.mark.parametrize('kind,native_exit', [('clean', 0), ('warn', 4)]) +def test_completed_fail_keeps_supported_policy_native_exit(evidence, kind, native_exit): + report, _ = evidence(kind) + data = json.loads(report.read_text()) + check = data['checks'][0] + entry = check['kernel_results'][0] + data['overall_verdict'] = check['verdict'] = entry['verdict'] = 'fail' + check['returncode'] = entry['returncode'] = native_exit + report.write_text(json.dumps(data)) + summary, code = helper.summarize(report) + assert code == 0 + assert '"entry_verdict": "fail"' in summary + + +@pytest.mark.parametrize('records,limit,expected_exit', [ + (3, 1, 2), (1, 1, 0), (1, 3, 0), (3, None, 0), (None, 1, 0), +]) +def test_report_only_checks_entry_record_limit(evidence, records, limit, expected_exit): + report, _ = evidence('warn') + data = json.loads(report.read_text()) + check = data['checks'][0] + entry = check['kernel_results'][0] + metadata = {'reported_count': 3, 'count_meaning': 'exact', + 'emitted_records': records, 'requested_limit': limit} + check['diagnostic_summary'] = entry['diagnostic_summary'] = metadata + report.write_text(json.dumps(data)) + summary, code = helper.summarize(report) + assert code == expected_exit + if expected_exit: + assert 'reported emitted records exceed requested limit' in summary + else: + assert 'Report consistency: INCONSISTENT' not in summary + assert '"entry_verdict": "warn"' in summary + assert '"finding_count": 1' in summary + + +def test_report_only_rejects_exact_aggregate_with_omitted_entry_truncation(evidence): + report, _ = evidence('warn') + data = json.loads(report.read_text()) + check = data['checks'][0] + entry = check['kernel_results'][0] + check['diagnostics_truncated'] = True + entry.pop('diagnostics_truncated') + metadata = {'reported_count': 1, 'count_meaning': 'exact', 'emitted_records': 1, + 'requested_limit': 64} + check['diagnostic_summary'] = entry['diagnostic_summary'] = metadata + report.write_text(json.dumps(data)) + with pytest.raises(ValueError, match='lower-bound marker contradicts exact count'): + helper.summarize(report) diff --git a/tests/instrumentation/rocjitsu_sanitizers/test_waitcheck_skill_recipe.py b/tests/instrumentation/rocjitsu_sanitizers/test_waitcheck_skill_recipe.py new file mode 100644 index 000000000..aa8331c43 --- /dev/null +++ b/tests/instrumentation/rocjitsu_sanitizers/test_waitcheck_skill_recipe.py @@ -0,0 +1,63 @@ +"""Validate the shipped draft recipe without running a checker or a kernel. + +Run: python -m pytest tests/instrumentation/rocjitsu_sanitizers/test_waitcheck_skill_recipe.py +The temporary object is inert text, not a qualification control. +""" + +from __future__ import annotations + +import hashlib +from pathlib import Path + +import pytest +import yaml + +from aorta.instrumentation.rocjitsu_sanitizers.recipe import load_sanitizer_recipe +from aorta.triage.recipe import RecipeSchemaError + +_REPO = Path(__file__).resolve().parents[3] +_TEMPLATE = ( + _REPO / ".agents/skills/rocm-sanitizer-triage/assets/waitcheck-exact-entry.yaml" +) + + +@pytest.mark.parametrize("resolved_index", [False, True]) +def test_unresolved_exact_entry_cannot_load(tmp_path: Path, resolved_index: bool) -> None: + data = yaml.safe_load(_TEMPLATE.read_text(encoding="utf-8")) + if resolved_index: + data["sanitizer_plan"]["source"]["kernel"]["code_object_index"] = 0 + recipe_path = tmp_path / "draft.yaml" + recipe_path.write_text(yaml.safe_dump(data), encoding="utf-8") + field = "entry_offset" if resolved_index else "code_object_index" + with pytest.raises(RecipeSchemaError, match=field): + load_sanitizer_recipe(recipe_path) + + +def test_resolved_template_preserves_exact_identity_and_bounds(tmp_path: Path) -> None: + # Only exercise recipe loading. These bytes/offsets are synthetic, not GPU code. + payload = b"inert recipe identity fixture" + (tmp_path / "fixture.co").write_bytes(payload) + digest = hashlib.sha256(payload).hexdigest() + data = yaml.safe_load(_TEMPLATE.read_text(encoding="utf-8")) + plan = data["sanitizer_plan"] + plan["target"] = "gfx1250" + plan["source"]["kernel"] = { + "name": "synthetic_entry", + "code_object": "fixture.co", + "code_object_sha256": digest, + "code_object_index": 1, + "entry_offset": "0x40", + } + recipe_path = tmp_path / "prepared.yaml" + recipe_path.write_text(yaml.safe_dump(data), encoding="utf-8") + recipe = load_sanitizer_recipe(recipe_path) + identity = recipe.kernel_specs[0] + assert identity.code_object == str((tmp_path / "fixture.co").resolve()) + assert identity.code_object_sha256 == digest + assert identity.code_object_index == 1 + assert identity.entry_offset == 0x40 + assert recipe.sanitizers == ("waitcheck",) + assert recipe.top_n == 1 + assert recipe.timeout_seconds == 120 + assert recipe.waitcheck_max_diagnostics == 64 + assert recipe.report_name == "sanitizer_report.json"