Fix for issue #111 - #112
Conversation
Modify example_stats.yml to include stat names separately and get_stats.py to parse the modified yaml. Signed-off-by: Rajarshi Das <drajarsh@gmail.com>
|
Hi, The build was finished. Unfortunately, the Jenkins log is not viewable public. You may ignore the failed message. You don't have to worry about it. We'll assign someone for a review and provide any feedback if needed. Thanks |
|
Thank you @JRPan |
Modify correl_mappings.py by removing regex symbols so that plot-correlation.py can use the csv output from the fixed get_stats.py script. Signed-off-by: Rajarshi Das <drajarsh@gmail.com>
|
@JRPan This change also allows completely removing the regular expressions from the correl_mappings.py file making it easier to read the simulation specific stat names. Thanks |
|
Hi sorry, we've been busy. Looks like the Jenkins is still failing. Are you sure the correl_mappings.py is working? There were no plots generated by Jenkins. This could happen when formatting changes and the python script was not able to find the regex string. |
|
Thank you for your comment. Thanks. |
|
@JRPan Hence, if the jenkins job is explicitly looking for regex, it will not find any. Thanks. |
|
@JRPan Thanks |
* dummpy h100 config
* updated hopper trace config
* supporting code to run hopper traces
* added some hopper opcodes
* Add TMA tracing support
* Watch out for WARPSYNC.COLLECTIVE instructions during instrumentation,
don't instrument the instructions between it and its target instruction
(both inclusive).
Resolves the issue from https://github.com/purdue-aalp/accel-sim-framework-private/issues/5
* Add TMA addr compression support and only process global addrs
* add TMA trace parsing
* Add TMA opcodes
* fix parser for TMA
* Make the WarpSyncWatchdog optional.
The Watchdog should completely go away once the release version of NVBit
fixes the bug of WARPSYNC.COLLECTIVE.
* tracer optionally ignores TMA instructions
* Tracer: more robust behavior when leave_kernel_launch makes it terminate on CUDA error
* Tracer can optionally tracer register values
* Add new trace file version to tracer. AccelSim trace parser not yet
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 ...]
}
* stub code for consuming the remainder of version 6 trace
* fixed a trace file formatting bug introduced by AI coding tools
* continue fixing formatting bug
* add tma apps
* fix a typo
* add H100 to launch control
* add bulk opcode and set PF as nop
* fix a deadlock due to EXIT and BAR inst
* add LLM tool ignore list
* add more bulk test args
* add option to run ncu set stats
* fix typo
* add reduction metrics count for UTMAREDG and UBLKREDG
* fix for tma simulation
* use guard to run with public NVBit
* Vllm update (#485)
* add vllm
* update nvbit
* Update install_nvbit.sh
* add regular load test to tma benchmark
* Modified h100 trace config (#41)
* adding H100 CIs
* apply format
* ncu flag change (#45)
* enable weekly on h100-test
* format codes
* Add TMA tracing support
* Vllm update (#485)
* add vllm
* update nvbit
* Update install_nvbit.sh
* apply format
* add register value parsing
* add support for tracing uniform register
* add fence async proxy and mbarrier support
* add tma mbar addr
* collect arrive count
* let tool read zero register to avoid potential issue
* fix deadlock and add trace debug message for mbarrier
* fix a bug in tracer tool
* add more tma app config
* fix a bug with UBLK op setup and reduce the iter count for tma ubench
* update tma app path and revert the iter change
* fix execution mask off issue
* handle additional variants of mbarrier arrive
* add run dir for monitoring run
* Fix a bug causing parsing regular register as uniform register
* set memory space for syncs op to be sharedmem
* tma: add cutlass test example
* tma: run gemm once only
* tma: update tracer with debug prints
* tma: pass the DYNAMIC_KERNEL_RANGE env to the script
* add additional SYNCS variants and opcodes
* add cluster information and multicast mask
* add parsing support for cluster information
* set multicast mask properly
* fix a bug in inject tool
* keep track of oob byte count
* pass oob count to the warp_inst_t
* add support for TMA bulk group
* add support for LDGSTS mbarrier based completion
* fix a bug in post-processing script
* fix a bug in tracer on kernel filtering
* update cultass TMA benchmark
* Reduce apps
* fix a bug in tracer
* Fix an issue causing assertion error
* add support for specifying additional ncu flags
* let set ncu to also use the flags
* fix a bug with parsing TMA address with no in-bound addresses
* Add GMMA modeling
* add gmma commit group support
* add additional build dir to ignore list
* Implement warpgroup.depbar, also double GMMA latency and halve their throughput to account for underestimating MMA warp duration
* Fix gmma initiation interval misconfiguration
* Fix a bug in tracer tool that parse TMA even if the instruction is predicated off
* add tracer runtime enable control
* Correctly implement inst_trace_t copy()
* format code
* add missing function for tracer in public nvbit
* address PR comments
* remove duplicated base_delta_compress_tma from rebase
* Remove custom CUTLASS kernels
* Enable nvbit by default with run_hw_trace
* On no kernel list file, write to stderr so it can be marked as failed
* enable instrumentation by default
* make gmma commit group depends on gmma opcode
* format tracer code
* Add CUPTI PM sampling tools (#57)
* adding H100 CIs
* adding Cesar's CUPTI timeseries profiler
* Add device ID support and PM sampling capability checks in CUPTI tools
* Add supported metrics for NVIDIA GeForce RTX 5090
* CUPTI include might have been moved
* Update CUPTI library search paths to include CUDAToolkit_LIBRARY_DIR
* adding hopper supported metrics
* Remove base metric from CUPTI injection settings and increase PM sampling buffer size
* 5090 metrics
* H200 metrics
* Add CUPTI PM sampling tools and documentation
- cupti.sh: wrapper script for easy profiling with env setup
- run_cupti.py: batch profiling for benchmark suites
- plot.py: visualization of PM sampling CSV output
- pm_injection.cu: fix flush timing (flush on kernel exit, not entry)
- pm_sampling.h: add ClearRanges() to prevent sample accumulation
- README.md: comprehensive usage docs and troubleshooting
Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>
* Update default metrics to cycles and instructions only
Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>
---------
Co-authored-by: Claude Opus 4.5 <noreply@anthropic.com>
* Add Hopper workflows and QV100/A100 support (#59)
* +x
* try again
* test
* rsync the folder
* refactor
* Enhance H200 trace generation script and update workflow to sync traces
* Fix PATH export in H200 trace generation script
* Update git clone URL in H200 trace generation script
* Update H200 trace generation script to use GPU_Microbenchmark
* Update H200 trace generation script to clone gpu-app-collection with submodules
* Add creation of dummy data_dirs in H200 trace generation script
* Move synced traces to trace depot in H200 workflow
* add rodinia to hopper workflows
* Clean up existing trace depot before moving new H200 traces
* try hopper weekly
* test if it work without actually running simulations
* try again
* checkout default branch
* weekly ready. enable everything back
* Add QV100 and A100
* Skip traces folders when walking hardware directory (#60)
* Skip traces folders when walking hardware directory
Traces folders are large and don't contain stats files needed for
correlation plotting. This prevents unnecessary traversal.
Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>
* Add forbidden files check to CI workflow
* Remove pull request trigger from CI workflow comments
* Update CI workflows to improve trace handling and job execution
* Optimize application build process by adding parallel execution for faster compilation
* Add 'cutlass' target to application build in H200 trace script
* Remove cutlass_100_hopper_tma_warpspecialized_gemm from small TMA configuration
* make cutlass_mini
* Refactor trace handling in hopper weekly workflow to streamline directory management and improve trace path usage
* rm device-*
* add tma uBench to CI
* Adding TMA tests to hopper-weekly CI workflow
---------
Co-authored-by: Claude Opus 4.5 <noreply@anthropic.com>
* Print kernel names to a file for cupti (#63)
Co-authored-by: Ni Kang <kang222@poto.ecn.purdue.edu>
* Fix tma tracing when instruction is predicate off (#65)
* Let tracer outputs correct format dummy values when TMA is predicate off or not using private NVBit
* add missing space
* add assertion for requiring SYNC instruction to use register value tracing
* format code
* Add fork owner's default branch as fallback in gpgpu-sim checkout
When the matching branch name is not found in the fork owner's
gpgpu-sim repo, try their default branch before falling back to
upstream dev.
Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>
---------
Co-authored-by: JRPAN <25518778+JRPan@users.noreply.github.com>
Co-authored-by: Claude Opus 4.5 <noreply@anthropic.com>
* print counters
* close the gz counter file at the end
* add plot generation script
* add normalize option to the plots
* Set TMA instruction based on opcode with public NVBit (#70)
* Set TMA instruction based on opcode with public NVBit
* format code
* Add value tracing for mbarrier.test_wait in SASS (#72)
* Add value tracing for mbarrier.test_wait in SASS
* Add more descriptive error msg and comments
* Address PR
* Remove redundant check for SYNCS instructions
* CI: Add correlation, stats archiving, and tracing improvements (#75)
* Update CI workflows to improve trace handling and job execution
* Optimize application build process by adding parallel execution for faster compilation
* Refactor trace handling in hopper weekly workflow to streamline directory management and improve trace path usage
* correlation in CI
* always enable cache control
* gmma is on a different build target
* enable reg value tracing
* Barrier fix (#77)
* Barrier fix
* Removed log files
* format
* Delete barrier_reaches.log
* Delete debug.log
* Delete gpgpu_inst_stats.txt
* Delete perf_counter.csv.gz
* Delete util/hw_stats/pm_cupti_tools/pm_metrics_h100.txt
* Removed duplicate comments
---------
Co-authored-by: Ni Kang <kang222@poto.ecn.purdue.edu>
Co-authored-by: JRPan <25518778+JRPan@users.noreply.github.com>
* Add LRC config knobs and additional hardware counter for LRC (#78)
* Add LRC knobs in config option file
* Add hardware specific metrics to collect
* Move LRC config knobs to a dedicated file and add plotting
* Bump makefile build standard to C++17
* Add normalization for LRC sector count and stall count
* H200 config, CMake fixes, and tracer cleanup (#74)
* Adding H200 config. Update CMake flags
* remove accidentally added script
* H200 for now
* CMake: Enable DEBUG definition for RelWithDebInfo build type
* Add ARRIVES.LDGSTSBAR.64.ARVCNT support in trace parser (#80)
* Add ARRIVES.LDGSTSBAR.64.ARVCNT support in trace parser
Handle ARVCNT opcode in the OP_ARRIVES case, setting both the
m_is_ldgsts_arrives_mbar flag and the new m_is_ldgsts_arrives_arvcnt
flag to distinguish arrive-on from complete-tx mbarrier operations.
Fixes: #73
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
* Refactor hardware data parsing to support multiple base directories
---------
Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com>
Co-authored-by: JRPAN <25518778+JRPan@users.noreply.github.com>
* Bar fix No.2 (#82)
* Barrier fix
* Removed log files
* format
* Delete barrier_reaches.log
* Delete debug.log
* Delete gpgpu_inst_stats.txt
* Delete perf_counter.csv.gz
* Delete util/hw_stats/pm_cupti_tools/pm_metrics_h100.txt
* Removed duplicate comments
* bar_fix no.2
* format
* update clang-format
* run docker job on pc0*
* Add assertion for bar_id in trace_driven.cc
Add assertion to validate bar_id range for named barriers.
* format code
---------
Co-authored-by: Ni Kang <kang222@poto.ecn.purdue.edu>
Co-authored-by: JRPan <25518778+JRPan@users.noreply.github.com>
Co-authored-by: Ni Kang <kang222@tgrogers-gpu01.ecn.purdue.edu>
Co-authored-by: WilliamMTK <an107@purdue.edu>
* PyTorch NVBit tracer hook and PM sampling improvements (#92)
* pytorch nvbit tool
* Add PyTorch NVBit Tracer Hook and example usage
* add wrapper
* Defer cuptiProfilerInitialize to context creation callback
cuptiProfilerInitialize fails with CUPTI_ERROR_UNKNOWN on CUDA 13.0
when called before any CUDA context exists. Remove the early call from
InitializeInjection() since initContext() -> sampler.SetUp() already
calls it when the context-created callback fires.
Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>
* Fix atexitHandler to avoid double-free by clearing context map without teardown
* Add debug logging for PM sampling context initialization and kernel launches
- Added a printf statement to log the initialization of the context for PM sampling, including the context pointer and device ID.
- Enhanced kernel launch callback to log the kernel name, device ID, current kernel count, and maximum kernels allowed in the session.
* block print until kernel is done
* debug prints
* Refactor branch checkout logic in CI workflow to simplify handling of default branches
* extend weekly to 15hr
---------
Co-authored-by: Claude Opus 4.5 <noreply@anthropic.com>
* Add IPOLY_modulo hashing knob (#98)
* CI improvements: forbidden file checks and hopper-weekly options (#102)
* Expand forbidden file patterns and add notebook size check
- Add patterns for binaries (.so, .a, .dylib), archives (.tgz, .tar.gz),
data files (.pdf, .ppm, .bin), and job outputs (.o[0-9]+, .e[0-9]+)
- Add debug_config/ and .trace$ to forbidden directories/extensions
- Add 512KB size limit check for .ipynb files to catch notebooks with
embedded Plotly/matplotlib outputs while allowing clean notebooks
- Organize patterns into logical groups with comments
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
* Add skip_tracer option and fix fork-aware branch selection in hopper-weekly
- Add workflow_dispatch input to skip tracer step (use existing traces)
- Update SASS-Hopper to run when Tracer-Hopper is skipped
- Fix fork-aware branch selection to use fork's default branch when branch not found
- Update success job condition to handle skipped jobs
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
* Fix fork-aware branch selection: fetch before checking branch exists
The ls-remote check was failing silently. Now we fetch first, then use
rev-parse to verify the branch exists locally. Also use 'remote show'
instead of 'ls-remote --symref' for getting the default branch.
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
---------
Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com>
* Add Claude Code rules for project documentation (#103)
Add structured documentation rules in .claude/rules/:
- accel-sim.md: Framework overview, build system, job launching
- gpgpu-sim.md: GPGPU-Sim source code organization and key classes
- trace-driven.md: Accel-Sim wrapper layer and trace replay
- tracer.md: NVBit tracer tool documentation
Rules are gitignored but force-added for sharing.
Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com>
* Multithread post-traces-processing with -j flag (#93)
* Multithread post-traces-processing with -j flag
Replace fork/exec/pipe stdin/stdout redirection with popen/pclose for
thread safety. Each kernel trace file is now processed in its own thread
with independent FILE* handles.
- Add -j N flag to limit parallel threads (default: 8)
- Use semaphore to control concurrency
- Add -pthread to Makefile
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
* Fix: exclude kernelslist.g when scanning directory
Skip output files (*.g) when finding kernelslist files to process,
preventing the tool from trying to re-process its own output.
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
* format
* Remove unused is_per_core parameter from post-traces-processing
The is_per_core variable was declared and parsed but never actually used
in any logic. This removes dead code and simplifies the CLI interface.
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
* Use thread pool instead of creating one thread per kernel
Reduces overhead when processing many kernels by reusing a fixed number
of worker threads instead of creating/destroying thousands of threads.
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
---------
Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com>
* Enable both replay-mode and cache-control for ncu profiling (#105)
Always use --replay-mode application --cache-control none for ncu
profiling to ensure consistent and accurate hardware counter collection.
Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com>
* Refactor SASS parsing and add SYNCS.ARRIVE.TRANS64.RED support (#106)
* Move SASS parsing to a separate function and add SYNCS.ARRIVE.TRANS64.RED
* Format code
* Add CUDA graph support and per-device output to PM sampling (#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 replay region support and CI improvements (#107)
* 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 (#108)
* 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>
* Add support for name and kernel-name-filter keys in application yaml file (#114)
* Add --per-kernel flag to run_simulations.py for parallel kernel execu… (#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>
* Fix CI build using wrong config due to inherited positional args (#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>
* Upgrade NVBit tracer to public v1.8 (#117)
* 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>
* Add AI-assisted PR review, nudge, and issue triage workflows (#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>
* ci: switch gpu-simulator build to cmake (#121)
* 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>
* Enhance issue triage skill documentation for clarity and detail
* Fix TMA instructions inheriting prior instruction's registers (#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) (#118)
* 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 (#6)
- Prefetch thread: wrap in try-catch to prevent std::terminate (#2)
- TracezReader: assign before open() to prevent leak on throw (#3)
- traceDsm: free ddict/dctx on early exit (#5)
- Add comments explaining 20x decompression fallback ratio (#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 #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>
* Comment out myocyte-rodinia-3.1 configuration in define-all-apps.yml
* Reduce trace memory pressure (#111)
* 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 #118, e1c2b678) 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 (#126)
* 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>
* Anchor L2 bandwidth regexes in plot-perf-counters (#131)
`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$`).
* SYNCS.EXCH.64 64-bit src2 plumbing and GMMA fixes (#129)
* 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 #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>
* Re-instrument when CUfunction or PC address differs (fix cuFFT LTO missed instrumentation) (#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 #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 #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>
* pr-ai-review: use shared AALP pr-review plugin (single-agent + stdout extraction) (#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>
* Fix flaky tracer sync and archive A100/V100 correlation plots in Hopper Weekly (#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>
* Promote spinlock detection to whole basic blocks (fixes #127) (#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 #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 #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 (#132)
* 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>
* Accel-Sim 2.0 release docs: README rewrite and citation metadata (#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 #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>
* Restore the pull_request trigger on the CI workflow
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>
---------
Co-authored-by: christindbose <chris241@purdue.edu>
Co-authored-by: Weili An <an107@purdue.edu>
Co-authored-by: Fangjia Shen <shen449@purdue.edu>
Co-authored-by: William-An <china_aisa@live.com>
Co-authored-by: Ahmad Alawneh <aalawneh@purdue.edu>
Co-authored-by: Connie120 <conniekang1023@gmail.com>
Co-authored-by: William-An <an107@rough-wolf-grows-fin-02.datacrunch.io>
Co-authored-by: Claude Opus 4.5 <noreply@anthropic.com>
Co-authored-by: Ni Kang <kang222@poto.ecn.purdue.edu>
Co-authored-by: Ni Kang <kang222@tgrogers-gpu01.ecn.purdue.edu>
Co-authored-by: yechen3 <liu2550@purdue.edu>


Fix for issue #111
The changes allow: