Accel-Sim 2.0: full NVIDIA Hopper (H100/H200) support - #553
Merged
Conversation
don't instrument the instructions between it and its target instruction (both inclusive). Resolves the issue from purdue-aalp/accel-sim-framework-private#5
The Watchdog should completely go away once the release version of NVBit fixes the bug of WARPSYNC.COLLECTIVE.
…ate on CUDA error
updated.
Trace file version will be 6 if envvar ALLOW_REG_VAL_TRACING is set to
one. In this version, new fields are appended after the last field from
version 5:
- Val|NoVal // whether this instruction trace contains register values
if dest_num > 0{
- 1|32 // Each thread has its own reg value. Coalesced to one if
all threads have same value.
- dest_reg_val [dest_reg_val ...]
}
for x in 0..(src_num-1){
- 1|32 // Each thread has own reg value. Coalesced to one if all
threads have same value.
- src_reg_x_val [src_reg_x_val ...]
}
…m#96) * Add CUDA graph support and per-device output to PM sampling - Handle CUDA graph stream capture: skip PM sampling during graph capture and sample during cuGraphLaunch instead - Support cuLaunchCooperativeKernel and cuLaunchKernelEx in addition to cuLaunchKernel - Per-device CSV and kernel name output files for multi-GPU/MPI workloads - Add LTS sector metrics to default sampling configuration - Increase default sampling interval from 300 to 3000 sysclk Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Update plot.py * Clean up metrics in cupti.sh Removed unused metrics from the CUPTI script. --------- Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com>
* Add mark_region spinlock handling mode to tracer Introduces a third spinlock handling mode that emits REPLAY_START/REPLAY_END markers to delimit spinlock regions in traces, enabling the simulator to replay these regions dynamically rather than fast-forwarding them entirely. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Refactor CI workflows into reproducible scripts Extract inline bash from workflows into standalone scripts that can be: - Run locally to reproduce CI flows - Called with subcommands for CI step visibility - Shared across workflows via lib/common.sh New scripts: - lib/common.sh: Shared functions (VPN, build, archive, correlation) - clone-gpgpusim.sh: Fork-aware branch selection (was duplicated 4x) - hopper-weekly.sh: Combined tracer + SASS job stages - main-sass.sh: Multi-GPU SASS simulation (QV100, A100, H200) - main-ptx.sh: PTX simulation with app building - main-tracer.sh: Tracer tool testing Usage: ./hopper-weekly.sh # run all stages ./hopper-weekly.sh sass # run SASS stages only ./hopper-weekly.sh sass-build # run single stage (for CI) Reduces hopper-weekly.yml from 246 to 90 lines Reduces main.yml from 416 to 130 lines Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Add simulator replay region support for spinlock simulation Introduces REPLAY_START/REPLAY_END pseudo-opcodes that enable the simulator to replay instruction regions when mbarrier acquisition fails. This complements the tracer's mark_region mode for simulating spinlock behavior. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Add H200 GPU mapping to correlation configuration * Update CI workflows to actions/checkout@v6 Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Add cancel-slurm-jobs.sh script for CI cleanup Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Add replay region support with deadlock detection Implement is_in_replay() override in trace_shd_warp_t and add iteration counter for replay loops with deadlock detection after 100 iterations. - Add m_replay_iterations counter and deadlock abort - Add N=224 entry to GMMA latency mapping table - Add debug prints for replay region enter/exit/loop - Remove trace_pc rollback in issue_warp (handled in scheduler now) - Remove unused vllm/offline.py Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Update hopper-weekly.sh to support H200 simulations * Treat EXIT as implicit REPLAY_END in replay regions Extract replay-exit logic into handle_replay_region_exit() helper to avoid duplication. When EXIT is encountered while in a replay region, apply the same replay logic before executing the EXIT instruction. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Fix EXIT handling in replay regions and add TMA warp detection - Re-enable deadlock detection with improved output (flush streams) - Fix EXIT in replay regions: force active threads to exit instead of replaying when EXIT is the last instruction - Add TMA warp detection by scanning for UTMALDG instructions - Remove verbose debug print statements Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Add debug symbols to Release builds for profiling * record used mbarriers. * Update CMake configuration for debug builds and remove TMA debug print * Add STAS instruction opcode and trace parsing for Hopper - Add OP_STAS to trace opcode enum - Add STAS to Hopper opcode map with STAS_OP category Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> --------- Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com>
* Add debug symbols for profiling and merge chiplet stats - Always include -g and -fno-omit-frame-pointer for profiling support - Merge CHIPLET_ACC stats into GLOBAL_ACC for hardware correlation - Add CHIPLET_ACC stat patterns to example_stats.yml Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Add PREEXIT/USETSHMSZ opcodes and fix TMA store destination - Add PREEXIT and USETSHMSZ opcodes to Hopper opcode map as NOPs - Fix OP_UTMASTG to clear destination register since TMA store has no register output Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Add REDG opcode and update GMMA latency mapping * Refactor barrier ID parsing logic and improve assert checks for trace format --------- Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com>
accel-sim#115) * Add --per-kernel flag to run_simulations.py for parallel kernel execution Enables launching each kernel as a separate simulation job by creating per-kernel subdirectories with individual kernelslist.g files. Also refactors job submission and trace directory lookup into reusable methods, and updates simulator binary caching to check mtime. Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com> * Add support for per-kernel directory processing in get_stats.py --------- Co-authored-by: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
…el-sim#112) In build_accelsim(), `source ./gpu-simulator/setup_environment.sh` without arguments inherits the function's positional parameters ($1="false" from use_srun). This causes setup_environment.sh to set ACCELSIM_CONFIG=false instead of release, building to bin/false/ instead of bin/release/. Explicitly pass "release" to prevent this. Co-authored-by: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
* Per-kernel stats, HMMA sizes, write bypass, and plot fixes
- Add per-kernel directory processing with parallel workers in get_stats.py
- Add more HMMA N sizes (24, 40, 80, 88, 104, 200) to hopper_opcode.h
- Bypass L1 cache for STRONG.GPU and BYPASS store ops in trace_driven.cc
- Fix plot-correlation.py CSV parser to rebuild klist per stat and handle all_kernels
- Reset accum_cycles to start from 0 in pm_cupti plot.py
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
* Add 2-hour time limit to CI gpu-lock srun jobs
Prevents CI HW stats/trace jobs from running indefinitely on the DGX node.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
* Upgrade NVBit tracer to public v1.8 release
Remove USE_PRIVATE_NVBIT guards and always enable TMA instrumentation
and cluster CTA functions, which are now available in the public NVBit
v1.8 release. Update install_nvbit.sh to download v1.8.
API changes from private NVBit:
- get_cluster_ctaid/get_cluster_nctaid: now in NVBit utils.h
- nvbit_add_call_arg_tma_param_handle -> nvbit_add_call_arg_tma_param_handle_and_size
- nvbit_parse_tma_{src,dst}_addrs: changed to callback-based API
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
* Use local /tmp workdir for simulations to reduce NFS pressure
Simulations now run in /tmp instead of directly on NFS. A background
sync loop copies results back to the NFS directory every 5 minutes.
On exit (normal, error, SIGTERM, SIGINT), a final sync runs and the
tmp dir is cleaned up. If the process is SIGKILL'd (e.g. OOM), the
sync loop detects the job is gone via squeue and handles cleanup.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
---------
Co-authored-by: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
…im#120) * Add AI-assisted PR review, review nudge, and issue triage workflows Three GitHub Actions workflows driven by the existing local opencode + Qwen3.5-27B setup: - pr-ai-review.yml: fires when a PR gets the `ready` label, runs a smoke-test review via opencode, posts the review as a comment, and either requests `aalp-grad` as reviewer (PASS) or adds `ai-concerns` (CONCERNS). Uses round-robin team assignment configured on the team. - pr-review-nudge.yml: scheduled Mon/Thu 14:00 UTC, pings requested reviewers on PRs labeled `awaiting-human-review` that have been idle for 3+ days, and cleans labels once a substantive review lands. - issue-triage.yml: on new issues, runs opencode with the triage skill and applies labels from a maintainer-defined allow-list. All API calls use `actions/github-script` with the default GITHUB_TOKEN; only the PR diff fetch uses curl (for the application/vnd.github.v3.diff accept header). Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * Run pr-review-nudge on tgrogers-raid self-hosted runner Keep all new workflows on the same runner for consistency. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * Check out base branch for AI review, not PR head The skill file lives on the base branch (ai-automation), not on the PR head. opencode only needs pr.diff + project conventions, so checking out the base is sufficient and lets the AI verify the PR against the current project rules rather than whatever stale rules the PR branch has. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * Checkout default branch for skills + add CONCERNS footer Two changes: 1. Always check out the repo default branch so skill files and project rules are present, regardless of PR head. Survives switching the default back to h100-test later. 2. When verdict is CONCERNS, append a footer to the PR comment telling the author how to proceed: either push a fix and re-apply `ready`, or reply dismissing the flag and manually request human review. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * Explicit default-branch checkout for issue triage Matches pr-ai-review.yml. Robust to the default branch changing later (e.g. switching back to h100-test after merging ai-automation). Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * Drop CLAUDE.md reference from skills CLAUDE.md is gitignored personal config, not in the repo. Project conventions live in .claude/rules/ instead. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * Use GitHub App token for team reviewer request Default GITHUB_TOKEN cannot resolve org teams; mint a short-lived installation token from the AALP GitHub App and use it only for the requestReviewers call. Comments and labels stay on GITHUB_TOKEN. App ID and PEM path are stored in repo secrets (AALP_APP_ID, AALP_APP_PEM_PATH) so they stay out of the workflow file when the repo eventually goes public. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * Drop --dangerously-skip-permissions from opencode calls Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> --------- Co-authored-by: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
* Add opencode CI analysis to hopper-weekly Appends AI-generated root-cause analysis (on failure) or simulation accuracy report (on success) to the Actions run summary. Skills live under .github/skills/ so they travel with the workflow. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * Add opencode CI analysis to main SASS-Simulation Same success/failure opencode steps as hopper-weekly, so push/PR runs get an AI summary on the Actions run page. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * Tee build output to build.log for opencode CI analysis GitHub Actions captures step stdout/stderr only on the GitHub side, so the opencode failure-analysis step (which can't use gh) had no way to see compiler/linker errors on a build failure. Tee the Build Accel-Sim step to build.log and point the skill at it. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * ci: switch gpu-simulator build to cmake The self-hosted runner's shared workspace was failing `make clean` with "Directory not empty" on ./build/release, blocking PTX-Simulation CI. Switch gpu-simulator build steps to the cmake flow (which CLAUDE.md already documents) and pre-clean via rm -rf to sidestep the flaky Makefile clean target. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * ci: drop --dangerously-skip-permissions from opencode invocations Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> --------- Co-authored-by: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
…chKernel uses CUkernel handles (accel-sim#521)
…sim#119) * Fix a bug that cause TMA instructions to have prior instruction's dst and src as its own registers * Remove legacy TMA opcode lookup map --------- Co-authored-by: JRPan <25518778+JRPan@users.noreply.github.com>
* Add per-warp zstd compressed trace format (.tracez) Post-processor outputs .tracez with per-warp sub-chunked zstd compression (default 2048 instructions per sub-chunk, configurable via TRACEZ_SUBCHUNK_INSTS env var). Text instruction format is preserved inside chunks, reusing parse_from_string() as-is. Simulator reads .tracez via WarpTraceStream with a sliding window that decompresses one sub-chunk at a time per warp using pread() on a shared fd (lock-free parallel warp reads). Old sub-chunks are discarded as the simulator advances, keeping memory proportional to subchunk_size * active warps instead of total instructions. Backward compatible: .traceg/.traceg.xz still supported. Post-processor defaults to .traceg output; pass --tracez for new format. Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com> * Add traceDsm for .tracez disassembly Two output modes: - Default: simulator-compatible .traceg text format - --annotate: human-readable format with labeled fields (PC, mask, active count, dst/src regs, opcode, memory info) Built as part of the tracer_tool Makefile. Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com> * Fix subchunk index mapping in WarpTraceStream The reader was computing sc_size_ via ceiling division of inst_count/num_subchunks, which doesn't match the post-processor's fixed subchunk_insts split. This caused assertion failures when accessing instructions near the end of a warp (e.g., TMA workloads). Fix: pass the actual subchunk_insts from the file header and use it directly as sc_size_. Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com> * Add double-buffer prefetch and replay-safe sub-chunking Post-processor: sub-chunk boundaries never split replay regions (REPLAY_START to REPLAY_END). If a replay region crosses the nominal 2048 boundary, the chunk extends to include REPLAY_END. Sub-chunk sizes are now variable and stored per-warp in the index footer. Simulator: WarpTraceStream uses double buffering — a background thread prefetches the next sub-chunk while the simulator consumes the current one. Each buffer has its own ZSTD_DCtx for thread safety. Sub-chunk lookup uses binary search over cumulative instruction counts. Index format change: per-warp entries now include per-subchunk instruction counts. Existing .tracez files must be re-post-processed. Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com> * Add -lzstd to Makefile for CI build The Makefile-based build (used by CI) was missing the zstd link flag. CMake builds already had it via trace-parser/CMakeLists.txt. Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com> * Address review feedback: error handling and cleanup fixes - pread_exact: replace assert with descriptive error message (accel-sim#6) - Prefetch thread: wrap in try-catch to prevent std::terminate (accel-sim#2) - TracezReader: assign before open() to prevent leak on throw (accel-sim#3) - traceDsm: free ddict/dctx on early exit (accel-sim#5) - Add comments explaining 20x decompression fallback ratio (accel-sim#13) Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com> * Remove dead code in TracezReader and traceDsm Remove unused members/accessors from TracezReader: filepath_, filepath(), subchunk_insts(), num_tbs(), reset_tb_counter(), data_start_. Remove unused <cstdlib> include from traceDsm. Replace decompression size fallback with abort (content size is always written by our compressor). Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com> * Move traceDsm to others/; drop .tracez magic footer - Relocate traceDsm from tracer_tool/ to others/traceDsm_tool/ with its own Makefile, matching the layout of other auxiliary tools (bbv_tool, spinlock_tool, etc.). Drop the traceDsm target from tracer_tool/Makefile. - Remove the TRACEZ_MAGIC footer sentinel. Format dispatch is already driven by the .tracez filename extension, so the magic was only acting as a truncation check. The footer is now 8 bytes (u64 index_offset) instead of 12. Existing .tracez files must be regenerated. - Replace unchecked fread calls in traceDsm.cpp with a fread_exact helper that aborts with a clear message on short read, fixing -Wunused-result warnings. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * Share .tracez format structs in tracez_format.h Extract warp_index_t / tb_index_t plus the index read/write helpers into a header-only tracez_format.h next to the post-processor (the producer of the format). The simulator reader and the traceDsm disassembler now include it via -I flags on the respective build files instead of redefining the structs locally. Addresses William-An's review comment on PR accel-sim#118 about the duplicated format definitions. * Add tracez_format.h include path to top-level gpu-simulator CMake The top-level accel-sim.out and accel_sim (python wrapper) targets transitively include warp_trace_stream.h but don't inherit the include path from the trace-parser / trace-driven subdirectories, because add_subdirectory() is called before include_directories() in the parent CMakeLists. Add the path at the top level too. Fixes CI build failure introduced by e70fb64e. --------- Co-authored-by: Claude Opus 4.6 (1M context) <noreply@anthropic.com> Co-authored-by: WilliamMTK <an107@purdue.edu>
* Reduce trace memory pressure: opcode interning, lazy reg_vals, and consumed trace release
Three optimizations to reduce memory usage during trace-driven simulation:
1. Opcode string interning: Replace per-instruction std::string opcode
with uint16_t opcode_id referencing a shared OpcodeInternTable
singleton. Eliminates millions of duplicate heap-allocated strings.
2. Fix tma_memadd_info leak + lazy reg_vals: Add missing delete of
tma_memadd_info in destructor. Wrap reg_dest_vals/reg_src_vals in
a unique_ptr<reg_vals_t> that is only allocated when register values
are present, saving ~48 bytes per non-SYNCS instruction.
3. Release consumed trace instructions: After consuming an instruction
in get_next_trace_inst(), move-assign an empty inst_trace_t to the
consumed slot (when not in a replay region) to free its heap
allocations as simulation progresses.
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
* Add --profile-perf and --profile-mem options to run_simulations.py
Add profiling support to the job launch infrastructure:
- --profile-perf: wraps simulation with perf record -g for function-
level CPU profiling (output: profile_perf.data)
- --profile-mem: wraps simulation with heaptrack for heap memory
profiling (output: profile_heaptrack.zst)
- --heaptrack-bin: path to heaptrack binary or AppImage for systems
without heaptrack in PATH
Each profiling mode appends a postfix (.perf or .mem) to the output
file names so both can coexist in the same sim_run directory.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
* Add compare_profiles.py for profiling analysis and visualization
New script that parses perf record and heaptrack profiling data from
simulation runs and generates comparison reports with plotly charts.
Supports two modes:
- Single-run: bar charts of absolute metrics per benchmark
- Comparison (-A vs -B): correlation scatter plots, grouped bar
charts, speedup chart, and per-function hotspot charts
Charts (ordered): scatter plots first, then bar charts, then
per-benchmark CPU hotspots (from perf report) and memory hotspots
(from heaptrack top allocation sites) for both runs.
Generates profile_report.html (full) and profile_summary.html
(scatter plots only) in the output directory.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
* Add simulation rate plots to correlation and perf counter scripts
plot-correlation.py:
- Add --sim-rate option that generates a per-kernel simulation rate
(inst/sec) horizontal bar chart, sorted slowest-first, to identify
which kernels are simulation bottlenecks.
plot-perf-counters.py:
- Add Simulation Rate (inst/sec) and (cycle/sec) time-series plots
derived from the new wall_clock_ms column in perf counter CSV.
Shows how fast the simulator runs over time — dips indicate
expensive program phases.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
* Support parallel perf+heaptrack profiling and rename CLI flags
run_simulations.py:
- When both --profile-perf and --profile-mem are specified, launch
two separate parallel jobs per benchmark (one for perf record,
one for heaptrack) instead of nesting them. Each job gets its own
.sim template (slurm.perf.sim, slurm.mem.sim) and justrun script
(justrun.perf.sh, justrun.mem.sh) to avoid overwrites.
- text_replace_torque_sim now accepts profile_mode parameter and
returns the template filename used for job submission.
compare_profiles.py:
- Rename -A to --baseline (-A kept as short alias)
- Rename -B to --other (-B kept as short alias)
- Matches --baseline-name and --other-name flag naming
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
* Include benchmark args in bar chart labels
Use benchmark/args as chart labels instead of just benchmark name,
so different input arguments for the same application appear as
separate bars rather than being clustered together.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
* Improve compare_profiles.py: parallelism, per-kernel stats, and fixes
Parallel parsing:
- Add -j/--jobs option with ThreadPoolExecutor for parallel benchmark
parsing (default: CPU count). Uses as_completed() for reliable tqdm
progress updates regardless of completion order.
Heaptrack fixes:
- Extract AppImage once at startup via resolve_heaptrack_print() to
avoid race conditions when multiple threads call
--appimage-extract-and-run simultaneously on the same temp directory.
- Remove QT_QPA_PLATFORM=offscreen which crashed the bundled Qt and
killed heaptrack_print before summary lines were printed. Just unset
DISPLAY instead.
Per-kernel simulation stats:
- parse_sim_stats() now returns {"app": {...}, "kernels": [...]} with
per-kernel differenced aggregate stats (cycles, insn, sim time).
- App-level rate correctly computed as total_insn/total_time (no longer
overwritten by last kernel's snapshot value).
- Add [Kernel] correlation scatter plot for per-kernel simulation rate.
Output file selection:
- Prefer unprofiled output files (exclude .perf.o* and .mem.o*) for
accurate simulation stats. Warn on stderr if only profiled outputs
are found.
Plot improvements:
- Scatter plots: markers only, full details in hover tooltip.
- Bar charts: indexed app labels for different args, full path in
hover tooltip. All titles prefixed with [App] or [Kernel].
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
* Make inst_trace_t move-only, remove manual copy constructor
Address review feedback: the manual copy constructor and copy
assignment operator required listing every field and would silently
miss new fields added to the struct.
Since inst_trace_t is never actually copied (only moved or
constructed in-place), make it move-only:
- Delete copy constructor and copy assignment operator
- Convert raw owning pointers (memadd_info, tma_memadd_info) to
unique_ptr, matching reg_vals — compiler now manages cleanup
- Remove manual destructor body (unique_ptr handles all deletes)
- Default move constructor and move assignment (already were default)
The compiler will now error if anyone tries to copy inst_trace_t,
making the maintenance hazard impossible.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
* Add simulation cycle correlation scatter plot
Add [App] Simulation Cycle Correlation scatter plot to comparison
mode. Points on the y=x line indicate identical cycle counts between
the two builds (correctness). Points off the line highlight
benchmarks with cycle count differences that need investigation.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
* Add log_scale option to scatter plots, use for cycle correlation
Add log_scale parameter to _add_correl_scatter(). When True, both
axes use logarithmic scale — useful for cycle correlation where
benchmarks span orders of magnitude (thousands to millions of cycles).
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
* Add geometric mean to speedup bar chart
Append a GeoMean bar (dark color) at the end of the speedup chart
showing the geometric mean of all per-app speedups. Also display
the value in the chart subtitle.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
* Dump summary.txt and summary.csv to output directory
Write two files alongside the HTML reports:
- summary.txt: human-readable table with per-app simulation rate,
speedup, cycle match, heap memory, and top CPU functions
- summary.csv: machine-readable CSV with all metrics for further
analysis (e.g., pandas, spreadsheets)
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
* Merge hotspot charts side-by-side and add absolute comparison
CPU hotspots:
- Merge baseline and optimized into one grouped horizontal bar chart
with functions as the union of both runs' top functions
- Add absolute samples chart showing raw sample counts side-by-side
(useful for seeing actual time reduction, not just % shift)
Memory hotspots:
- Same pattern: merged peak MB chart + absolute allocation calls chart
Functions are sorted by max overhead/peak across both runs so the
most important ones appear at the top.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
* Put baseline first (top) in grouped horizontal bar hotspot charts
In plotly horizontal grouped bars, the last trace renders on top
of each cluster. Swap trace order so baseline (last added) appears
above the other run's bar, making visual comparison more intuitive.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
* Add tqdm to requirements.txt
compare_profiles.py imports tqdm for progress bars during parallel
benchmark parsing. Without this, users get ModuleNotFoundError.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
* Fix _submit_job to use mode-specific .sim template for profiling
After the rebase, _submit_job() (from upstream's per-kernel feature)
hardcoded the global job_template variable. When --profile-perf or
--profile-mem are used, text_replace_torque_sim() writes mode-specific
templates (slurm.perf.sim, slurm.mem.sim) but _submit_job() was
always submitting the unmodified slurm.sim instead.
Fix: add sim_template parameter to _submit_job(), pass the return
value of text_replace_torque_sim() through both call sites (per-kernel
and standard mode). Falls back to global job_template when None.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
* Remove opcode interning and lazy reg_vals from trace-driven simulation
Reverts the per-instruction memory optimizations introduced in b9383ac2
"Reduce trace memory pressure":
1. OpcodeInternTable / uint16_t opcode_id -> std::string opcode
2. unique_ptr<reg_vals_t> -> direct reg_dest_vals /
reg_src_vals vectors
3. Move-empty over consumed warp_traces[trace_pc] in
trace_shd_warp_t::get_next_trace_inst() -> removed
Those optimizations targeted the high-water-mark of the old design that
loaded an entire warp's trace into memory upfront. The per-warp .tracez
paging system (PR accel-sim#118, e1c2b67) keeps only subchunk_size x active_warps
instructions resident at any time, so the per-instruction byte savings
are no longer meaningful while the indirection (get_opcode(),
has_reg_vals(), get_reg_dest_vals(), get_reg_src_vals(), the global
mutable singleton) imposes ongoing readability cost.
Preserved:
- Move-only inst_trace_t and unique_ptr<inst_memadd_info_t> /
unique_ptr<tma_inst_memaddr_info_t> ownership from a6c786b9 -- these
are general code-health wins independent of paging, and the
unique_ptr-managed tma_memadd_info also subsumes the leak fix from
b9383ac2 part 2.
- Tracer-side opcode_id under util/tracer_nvbit/ -- unrelated NVBit
runtime instrumentation ID.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
---------
Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com>
Co-authored-by: JRPan <25518778+JRPan@users.noreply.github.com>
* Remove scan_warp_metadata helper scan_warp_metadata was only invoked on the text trace path and was disabled (commented out) on the .tracez path to avoid decompressing all sub-chunks at init time. The metadata it produced (is_tma_warp flag and kernel_info_t::register_used_mbarrier_addr) is no longer consumed after the corresponding cleanups in gpgpu-sim, so the helper and its TODO are removed. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * Drop unused-mbarrier TRYWAITs in post-traces-processing After the simulator-side cleanup, SYNCS.PHASECHK.*.TRYWAIT instructions targeting mbarriers that nothing ever ARRIVE-s on will deadlock the simulator (no phase advance). Detect and strip them at trace post- processing time so the resulting .tracez/.traceg is safe to replay: - Per kernel file, scan all warps' instructions and collect the set of mbarrier addresses (low 32 bits) that any SYNCS.ARRIVE*, ARRIVES.LDGSTSBAR, or UTMALDG (TMA hardware arrive on tma_mbar_addr) targets. - In a second pass, drop any TRYWAIT whose mbarrier addr is not in that set. Print one stdout line per dropped mbarrier, prefixed with the kernel trace filepath so output stays attributable when threads run concurrently. - Initialize the previously uninitialized `lineinfo` local to 0. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> --------- Co-authored-by: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
`df.filter(regex=...)` uses `re.search`, so the unanchored patterns `partiton_replys_in_parallel` and `partiton_reqs_in_parallel` matched the companion counters registered alongside them in gpu-sim.cc — the `_total` cumulative variants, plus `_util` / `_util_total` for the reqs case. The plotted L2 bandwidth was being inflated 2x (replies) and 4x (reqs), and mixing per-cycle and cumulative totals made the .diff() output meaningless. Anchor both patterns with `^…$` so each filter selects exactly the one intended counter, matching the convention already used elsewhere in this file (e.g. `gpu_sim_cycle$`, `GLOBAL_ACC_R_MISS$`).
* Enhance issue triage skill documentation for clarity and detail * tracer: capture 64-bit src2 for SYNCS.EXCH.64 via generic aux field Add auxRegVals[32] to the regular trace record as an opcode-defined per-lane slot. SYNCS.EXCH.64 populates it with Rn+1 of its src2 pair; the writer fuses srcRegVals[1] (lo) and auxRegVals (hi) into one 64-bit column. Also enable reg-val tracing by default. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * Widen reg_val_t to uint64_t for SYNCS.EXCH.64 64-bit src2 Widen reg_val_t to std::array<uint64_t, WARP_SIZE> so the parser can carry the full 64-bit src2 value for SYNCS.EXCH.64 (and future 64-bit operand slots). Narrow explicitly at consumption sites that store into 32-bit operand fields (init.count, arrive.count/txCount, complete_tx.txCount, wait.phase). Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * Drop intermediate thread_counts copy in SYNCS.EXCH.64 reversal Read trace.reg_src_vals[1][i] directly instead of copying into a std::array<uint64_t, WARP_SIZE> local first. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * Capture mbarrier init_as_one flag from SYNCS.EXCH.64 trace Per issue accel-sim#123, the captured 64-bit register holding the mbarrier init state encodes the pending-thread-count's top bit (bit 63) as a phase- advance trigger. Extract that bit into syncs_operand.init_as_one for the consumer in shader.cc to honor. Also assert the lower 20 bits of the pending-thread-count field mirror the expected-thread-count field, which is the only init encoding we have observed. A failure here is not necessarily a bug — it just means hardware produced a pattern we have not modeled and the decode below needs to be extended. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * Handle large GMMA N sizes and gate SYNCS.EXCH.64 assert on lane mask Compute latency/initiation interval as N/2 for GMMA N>=64 instead of failing the lookup table. Skip the SYNCS.EXCH.64 pending-thread-count sanity check for inactive lanes whose register values are undefined. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * traceDsm: write .traceg next to the input file by default MODE_SIM now derives the output path from the input (.tracez -> .traceg) instead of dumping to stdout, and rejects inputs without a .tracez suffix. --annotate still goes to stdout. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * CI: drop redundant lastSuccess symlink loop and tighten failure skill The post-correlate stage in main-sass.sh used to mv the workspace into lastSuccess; that left $PWD empty so a follow-up loop symlinked the files back. After the move was switched to rsync the workspace is no longer emptied, but the symlink loop stayed and now fails with "File exists" on the first item. Also update ci-failure.md so the AI analyzer anchors on the literal exit-1 error, knows about main-sass.sh / main-tracer.sh / h200.sh, and checks whether sims actually succeeded before blaming simulator code. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> --------- Co-authored-by: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
…ssed instrumentation) (accel-sim#134) * Re-instrument when CUfunction or PC address differs NVBit tracers cached instrumented kernels in a set keyed only on the CUfunction handle. With cuFFT LTO workloads, the driver can reuse a CUfunction handle (or repurpose a kernel's binary address) across launches, so a handle already in the set would skip instrumentation even though the underlying SASS had been rewritten. This produced intermittent failures: CUDA_ERROR_UNKNOWN on launch, or kernels with 0 recorded instructions. Key the already_instrumented set on (CUfunction, PC address) pairs so a kernel is re-instrumented whenever either the handle or its function address differs. Add a shared accelsimUtils::PairHash in a new util/tracer_nvbit/include/utils.h and adopt it in both tracer_tool and spinlock_tool, with the include path wired up in their Makefiles. Refs accel-sim#133 Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * Include mangled name in instrumentation dedup key Address review feedback: a (CUfunction, PC address) pair can still collide when a program reuses the same handle and address for a different kernel. Add the mangled kernel name (nvbit_get_func_name with mangle=true) as a third key component so such cases are treated as distinct functions and re-instrumented. Switch already_instrumented to std::tuple<CUfunction, uint64_t, std::string> and replace PairHash with a generic accelsimUtils::TupleHash (boost-style hash combine) in the shared util/tracer_nvbit/include/utils.h. Refs accel-sim#133 Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> --------- Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
… extraction) (accel-sim#136) The review skill + reviewer agent now come from purdue-aalp/AALP-claude-plugins, kept checked out on the runner at $HOME/AALP-claude-plugins (override with the AALP_PLUGINS_DIR repo variable) — one shared source of truth instead of a copy in each repo/branch. The skill and agent are copied into the workspace before the run (opencode refuses to read outside --dir headless). opencode runs a single read-only pr-review agent that PRINTS the review ending in a VERDICT line; the workflow extracts pr-review.md / pr-verdict.txt from that output deterministically. The model never calls a write tool, so it cannot 'forget' to — the root cause of the 'AI reviewer did not produce output' failures. Verified end-to-end against the local model. Removes the now-unused .github/skills/pr-review.md. (An orchestrator/reviewer subagent split was tried first but opencode's task tool hangs with the local llama.cpp provider, so it was dropped.) Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…er Weekly (accel-sim#137) * Harden tracer remote_sync against flaky SOCKS tunnel The Hopper Weekly tracer-sync-push step has failed for four consecutive weeks with broken-pipe / socket-IO errors while rsyncing the repo to the H200 box over the userspace SOCKS5 proxy. The push included the 313 MB .git history, which the remote never uses (h200.sh only builds the tracer and clones gpu-app-collection itself), so a large unneeded transfer was choking the fragile tunnel with no retry. Exclude .git from the transfer, keep partial files and time out stalled sockets (--partial --timeout=120), and retry up to 5 times with backoff so a transient tunnel drop resumes instead of failing the weekly run. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * Archive A100 and V100 correlation plots in Hopper Weekly The weekly SASS job already simulates QV100 and A100 regressions alongside H200, but only collected stats and correlation plots for hopper-h100. The ampere-a100 and v100 slots in the statistics-archive were therefore going stale. Factor the per-GPU stats collection into archive_config_stats() and run it for hopper-h100, ampere-a100, and v100. Correlate all three (each guarded so one GPU's failure does not drop the others or block the push) and push their combined_per_kernel / combined_per_app plots to the archive. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> --------- Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
accel-sim#128) * Promote spinlock detection to whole basic blocks Spinlock detection only flagged instructions whose summed counts strictly differed across two runs, so any loop-body instruction landing on the same total (e.g. uniform back-edge predicates) slipped through. Downstream fast-forward in tracer_tool resets per-warp counters on any unflagged instruction, so a single missed instruction inside the loop body disabled fast-forward entirely (issue accel-sim#127). Capture each kernel function's CFG via nvbit_get_CFG at instrumentation time, persist it as a .bbs sidecar next to the histogram, and in spinlock_check promote the count-diff set to the union of all instructions in any basic block containing at least one diff'd instruction. Falls back to per-instruction emission for kernels with degenerate CFGs (BRX/JMX). Output file format is unchanged, so the tracer_tool consumer needs no changes. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * Add static SASS-pattern matcher for spinlock detection Introduces a domain-agnostic BBPatternMatcher library (bb_pattern_matcher.h/cpp) for flagging NVBit basic blocks by registered named predicates, and wires it into the spinlock tool with one shipped pattern (Hopper NANOSLEEP+PHASECHK+BRA wait loop). Pattern hits are persisted in the .bbs sidecar, unioned into the final spinlock_instructions.txt, and a warning is printed (and mirrored into the output file as '#'-prefixed comment lines) when the static matcher catches a BB that count-diff missed. The tracer_tool reader now skips '#' comment lines. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * Address PR accel-sim#128 Copilot review comments - Use %zu for size_t prints in spinlock_check (was %d, undefined on 64-bit). - KernelBBLayout::deserialize now returns bool; fails when the header line doesn't match, so a malformed .bbs no longer gets keyed under the default name "dummy" and silently disables BB promotion. - Body lines are dispatched by regex prefix instead of consuming line 2 as Degenerate: positionally — a missing or out-of-order Degenerate line no longer eats the first BB. - Reword degenerate-CFG warning from "kernel %s" to "function %s" since instrument_function_if_needed walks both kernels and device callees. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * format code * Add missing include --------- Co-authored-by: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
* Add -m metric group selection to cupti.sh Adds a -m flag (repeatable) to select among gpc/fbp/hub/dram/nvlink metric groups instead of hard-coding the metrics list. Defaults to nvlink when no -m is provided. Bumps PM_SAMPLING_MAX_SAMPLES to 160000 and INJECTION_KERNEL_COUNT to 20. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * Default -m to nvlink when omitted Without -m, the script previously fell into a no-op "default" group, contradicting the help text that promised nvlink as the default. Make the default actually nvlink and drop the unreachable case. Also adds sm__pipe_tensor_cycles_active_realtime.sum to the base metrics. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * No-op default for missing -m; print help instead Reverse the prior auto-default to nvlink. When -m is omitted, print the usage blurb and run with the base metrics only (cycles, insts, tensor). Factor the help text into print_usage so usage() can still exit on error paths. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> --------- Co-authored-by: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
…el-sim#145) * Accel-Sim 2.0 docs: release notes, README rewrite, mark_region default - release.notes.md: add v2.0.0 entry with Table 1 and categorized changelog - README: reorganize for the Hopper release — Quick Start, full Hopper support callout (TMA/WGMMA/mbarrier), a dedicated "Tracing LLMs with vLLM" guide, GPUVision + compressed-trace guides, cumulative How-to-Cite section, roadmap (experimental Blackwell + multi-GPU incoming) - run_hw_trace.py: default --spinlock_handling to mark_region - torch_hook/run.sh: generalize to run any command; SPINLOCK_HANDLING_MODE=2 (mark_region), ALLOW_REG_VAL_TRACING=1, drop dead env var - embed local AccelWattch SVG and an example PM-sampling PNG (GitHub strips JS) Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * Add citation metadata files and restructure the README How to Cite section Add CITATION.cff, which drives GitHub's "Cite this repository" button, and CITATION.bib carrying the full cumulative set (the button only exports the preferred citation). Credit the software to an entity rather than an enumerated author list so it does not go stale. Split How to Cite into a base section and a power-model section, each with its own BibTeX block, replacing the table and ASCII diagram that duplicated each other. Order entries newest first. Fix the Accel-Sim 2.0 entry: it was an @inproceedings with booktitle "arxiv.org" and comma-separated authors, which BibTeX parses as a single author. It is now an @misc with the arXiv eprint 2608.22602, verified against the arXiv API. Also add the missing comma before the AccelWattch DOI, which was a parse error. Document the run.sh and cupti.sh wrappers in Quick Start and GPUVision, including the environment variables each sets and their defaults. Replace the ISCA 2020 PDF link, which returns 401, with the open copy that accel-sim.github.io itself links to. Reflow prose to one sentence per line so edits produce proportional diffs. Rendered output is unchanged. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> * Update the GPUVision docs for the new cupti.sh metric groups cupti.sh gained a repeatable -m flag (gpc/fbp/hub/dram/nvlink) in accel-sim#132, and its sampling defaults moved: INJECTION_KERNEL_COUNT 10 -> 20 and PM_SAMPLING_MAX_SAMPLES 80000 -> 160000, with a tensor-pipe metric added to the base set. Document the groups and correct the table. Also correct the override guidance. The five PM_/INJECTION_ settings are unconditional exports, so passing them in the environment has no effect -- the script overwrites them. Only PM_SAMPLING_CSV_PATH, which stays commented out, can be set that way. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> * Move the PM sampling example image into docs/img Keeps documentation images together with the other two the READMEs use (h100_correlation.png, accelwattch-flowchart.svg) instead of leaving a 266 KB screenshot sitting in the tool source directory. Git records this as a rename, and the reference in the pm_cupti_tools README is repointed at the new location. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> --------- Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
The trigger was disabled while this branch lived only in the private fork, where every branch was same-repo and `push` covered all of it. That no longer holds: a PR opened from a fork does not fire `push` in the base repo, so the release PRs into accel-sim/accel-sim-framework would run no CI at all. The workflow already handles the event -- it branches on github.event_name == 'pull_request' when choosing the report URL and when sending success/failure mail -- so only the trigger was missing. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
JRPan
added a commit
that referenced
this pull request
Aug 26, 2026
The 2.0 rewrite (#553) made slurm.sim hard-bound to a real Slurm + NFS cluster, which breaks every `run_simulations.py -l local` run. The local launcher submits through procman.py but still uses slurm.sim as its job template, and procman provides neither $SLURM_JOB_ID, squeue, nor a guarantee that rsync exists: * `#SBATCH --output=/dev/null` / `--error=/dev/null` -- procman parses exactly these lines to choose where to write a job's stdout/stderr, so all simulator output was being discarded. * The background watchdog `while squeue -j $SLURM_JOB_ID; ...` exits immediately when squeue is absent (e.g. inside a CI container) and falls straight through to `sync_to_nfs && rm -rf "$TMP_DIR"`, deleting the working directory while the job is still starting up. * procman rewrites the bare `$SLURM_JOB_ID` but not the braced `${SLURM_JOB_ID}` on the TMP_DIR line, so the staging directory and the log file inside it disagree about the job id. The net effect is that no `<name>.o<jobid>` ever lands in the run directory and job_status.py reports NOT_RUNNING_NO_OUTPUT for every job. This is what is currently breaking the gpgpu-sim CI, which clones this repo unpinned and launches with `-l local`. The old template only needs coreutils `mv`, so it works under both sbatch and procman. Reverting restores the local launcher; the tmp-staging behaviour can be reintroduced guarded on `[ -n "$SLURM_JOB_ID" ]` plus `command -v squeue rsync`. Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
7 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
This is the Accel-Sim 2.0 line from
purdue-aalp/accel-sim-framework-public:h100-test, proposed for merge intodev. It is large (146 commits) because it is the accumulated Hopper work rather than a single feature.Paper: https://arxiv.org/abs/2608.22602
Requires a matching GPGPU-Sim change
This depends on
h100-testin the GPGPU-Sim fork — the performance model half of the same work lives there, so this PR does not build against currentgpgpu-sim_distributiondevon its own:Both branches are synced with their respective upstream
devas of this PR.What lands here
Hopper execution model. TMA (bulk tensor movement, bulk/store groups, CGA shared-memory multicast), async WGMMA with warpgroup commit/wait and N-dependent MMA latency,
mbarrierproducer/consumer synchronization with async-proxy fences and remote arrive, threadblock clusters with cluster-aware CTA scheduling and distributed shared memory.Memory subsystem. HBM3/HBM3e timing, the L2 Request Coalescer, IPOLY+MODULO L2 hashing, and a chiplet / uGPU partitioned L2.
Tracer. Rebuilt: per-warp zstd-compressed
.traceztraces with page-by-page loading so simulator memory stays bounded regardless of kernel size, plustraceDsmfor decoding. Spinloop handling formbarrierpolling — detected regions are marked and the simulator re-evaluates the wait dynamically, rather than baking an unrepresentative poll count into the trace.LLM / PyTorch tracing. NVBit attaches to a live PyTorch process and traces only selected layers, so vLLM inference and full training steps can be traced without a standalone CUDA binary.
GPUVision. Cycle-level CUPTI PM sampling for time-series (not per-kernel-aggregate) correlation against the simulator.
Validation. 34,000+ kernel instances across 22 benchmark suites against real H100 silicon.
Review notes