Accel-Sim 2.0: Hopper (H100/H200) performance model - #143
Merged
Conversation
…2_rop_latency, dram_latency, gpgpu_dram_timing_opt, Gpgpu_num_reg_banks
* Updated h100 config * Update DRAM bus width and burst length in config Manual set to HBM3's config as we don't currently have HBM3 in tuner --------- Co-authored-by: root <root@quick-star-grows-fin-02.datacrunch.io> Co-authored-by: WilliamMTK <China_Aisa@live.com>
* Add LRC implementation to coalesce L2 requests * format code * mark merged request as reply when sending back * LRC applies to read request only * add dram counter to simple dram model * Add new statistics class * Add LRC stats with new statistics classes * Fix type conversion issue * Move new statistics counter to upper level and add a single non-templated counter class * Add new Statistics counter support for perf counter sampling * Add PerfCounter dump support for MultiUnitStatsCounter * Format code and remove unneeded type * Remove modulo operation * Add new average entry size counter for LRC * Address PR comments * Put to_string() in top class * Use C++17 for makefile build * Move handle_mf out as a lambda * Move LRC reply handling logic to a dedicated function
* create base class to operand collector * Add detailed and simple operand collector models to shader core configuration * Refactor memory access queue from list to vector and simplify iterator usage * Replace std::list with std::deque for response FIFO in ldst_unit and simt_core_cluster * Enhance pipelined SIMD unit with latency tracking and pipeline management * Enhance pipelined SIMD unit cycle logic to check result port availability * Refactor data structures in cache and scoreboard to use unordered sets and maps for improved performance * Refactor interconnect stall tracking and global writeback pipeline Rename stall counters to reflect directionality (icnt2mem, mem2icnt, icnt2core, core2icnt), move icnt_cycle to ICNT clock domain and mem-to-icnt pop to L2 clock domain, add round-robin fairness to Perfect_Advance in local_interconnect, refactor m_next_global from single pointer to deque to allow draining multiple global writebacks per cycle, and extract writeback_complete() helper to deduplicate pending-writes/scoreboard/LDGSTS completion logic. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Change m_pending_ldgsts from unordered_map to map for instruction count tracking * L2 config tuning for correct L2 BW * running formatter * remove occupied completely * format --------- Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com>
…ion (accel-sim#11) * Add ARRIVES.LDGSTSBAR.64.ARVCNT support for mbarrier arrive-on operation ARVCNT decrements the mbarrier pending_thread_count (arrive-on), unlike TRANSCNT which decrements tx_count (complete-tx). This corresponds to the cp.async.mbarrier.arrive.noinc PTX instruction. Fixes: purdue-aalp/accel-sim-framework-private#73 Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Run clang-format on shader.cc Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Refactor CI workflow and build scripts for improved clarity and functionality --------- Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com> Co-authored-by: JRPAN <25518778+JRPan@users.noreply.github.com>
The icnt_to_lrc sector counter did not account for TMA multicast requests being delivered to multiple CTAs. Multiply the sector count by the number of destinations in the multicast CTA mask so the counter matches the hardware behavior of lrc__xbar2gpc_sectors_op_read. Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com> Co-authored-by: JRPan <25518778+JRPan@users.noreply.github.com>
* Barrier fix * bar_fix no.2 * format * Format get_shader method for improved readability * format again * update clang-format. Enforce verison. --------- Co-authored-by: Ni Kang <kang222@poto.ecn.purdue.edu> Co-authored-by: Ni Kang <kang222@tgrogers-gpu01.ecn.purdue.edu> Co-authored-by: JRPAN <25518778+JRPan@users.noreply.github.com>
…sim#16) * Update LRC configs and use multimap + mf uid check for LRC lookup * Handle resend and edge cases * Use templated function to allow inline lambda * Remove redundant check for mbarrier existence * Fix typo in impl functions * Use kernel-level cluster mbarrier lookup table for fast mbarrier lookup * format code * Add additional check and avoid double completion on mbarrier * Bump LRC queue size to 1024
* Add IPOLY_MODULO hashing for L2 subpartition id hashing * Let H100 and H200 uses IPOLY_Modulo * format code
…sim#15) * chiplet read support * writes seems to be working * Add latency-modeled inter-chiplet queues with configurable size - Introduce LatencyQueue class for FIFO with cycle-based latency modeling - Replace simple deques with LatencyQueue for inter-chiplet communication - Add config options: inter_chiplet_queue_size, inter_chiplet_queue_latency - Add statistics counters for inter-chiplet read/write requests - Fix flow control: check queue space before pushing replies/writes - Remove CTA finish debug printf Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * fixed deadlock, but with 4 queues. Should be working with 2 queues. * 4 queues are required after all * clear memreqaddr before set * Add crash handler for debugging with backtrace on SIGSEGV/SIGABRT * format * Fix code review issues for chiplet support - Remove unused mf_handled variable - Fix typo: m_dest_chipet -> m_dest_chiplet - Fix FIXME: use actual cycle for status timestamps in L2interface/Chipletinterface - Fix incorrect description for -chiplet_mem_stride option - Make l2_cache chiplet members private - Handle full queue gracefully in forward_write_to_peer_chiplet Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * fix typo * Refactor chiplet queues into chiplet_icnt class Move the four chiplet-to-chiplet communication queues into a dedicated chiplet_icnt class for better encapsulation and cleaner initialization. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Add chiplet queue full and write fail stats Add per-subpartition counters for: - chiplet_queue_full: cycles when chiplet request queue was full - chiplet_write_fail: chiplet write forward failures Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Add src_chiplet to mem_fetch and refactor chiplet mapping - Refactor fold_to_chiplet into a dedicated function to avoid duplicating chiplet mapping logic across address decoding schemes - Add m_src_chiplet to mem_fetch, computed at construction time, to avoid recomputing chiplet ID on every access - Fix writeback mem_fetch allocations to pass correct sid/tpc instead of -1, ensuring proper chiplet routing for evicted cache lines - Re-enable cache reservation fail stats perf counters - Move chiplet config (n_chiplet, chiplet_interleave) into memory_config for cleaner access from mem_fetch constructor Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Add chiplet access tracking and fix GPC-based chiplet mapping - Add CHIPLET_ACC_R/CHIPLET_ACC_W memory access types to track cross-chiplet L2 accesses separately from local GLOBAL_ACC_R/W - Fix chiplet interleave mapping to use GPC grouping (tpc/8 % n_chiplet) instead of simple round-robin (tpc % n_chiplet) - Add L2_dram_queue_full statistics counter to track DRAM queue stalls - Move m_chiplet_id from l2_cache to baseline_cache for access in data_cache - Fix alternating priority logic to use static variable instead of cycle count Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Refactor chiplet address mapping: add get_dest_chiplet helper - Add get_dest_chiplet() method to memory_config for consistent chiplet lookup - Move chiplet calculation logic from mem_fetch constructor to config helper - Add fold_to_chiplet call in addrdec_tlx for proper TPC-based chiplet folding Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Fix L2 cache cycle priority alternation to be deterministic Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Revert CHIPLET_ACC_R/W and skip chiplet calc for single-chiplet - Remove the separate chiplet access tracking (CHIPLET_ACC_R/W) that was added to distinguish cross-chiplet L2 accesses - Skip chiplet src/dest calculations when n_chiplet == 1 to avoid unnecessary overhead in single-chiplet configurations - The GPC-based chiplet mapping fix is retained Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com> * Fix sub_partition calculation skipped when n_chiplet == 1 Move sub_partition calculation back outside fold_to_chiplet() so it always runs, matching the original behavior before the refactor in 307f0558. The early return in fold_to_chiplet was causing sub_partition to never be set for single-chiplet configs. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Fix use of uninitialized m_mem_config in mem_fetch constructor Use the config parameter instead of m_mem_config member, which is not assigned until later in the constructor. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Fix sub_partition for early-return modes in multi-chiplet configs For modes like IPOLY_MODULO that return early after setting sub_partition, recalculate sub_partition inside fold_to_chiplet after chip is modified. For n_chiplet == 1, early return preserves the mode's original sub_partition. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Update H100/H200 configs with chiplet and timing settings - Add chiplet partition settings (n_chiplet=2, queue_size=256, interleave=1) - Update clock domains from 1455 to 1980 MHz core clock - Change scheduler from gto to lrr with opndcoll_model=1 - Reduce kernel launch latency from 5000 to 3000 cycles - Increase L2 cache associativity from 192 to 384 - Add assertion in l2cache.cc to enforce 2-chiplet assumption Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Address PR review comments: ownership clarity and magic number - LatencyQueue: change sim_cycle/tot_sim_cycle from pointers to const references, add comment documenting non-owning reference semantics - addrdec.cc: replace magic number 8 with named constant sm_per_gpc, add comment explaining H100 has 8 SMs per GPC (GPC not yet modeled) Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> --------- Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com>
…im#19) mbarrier (SYNCS/ARRIVES) and fence instructions routed through the ldst_unit may not have output registers, causing writeback_complete() to never mark them as completed. This led to these instructions stalling in the pipeline. Add an else-if branch (analogous to the existing LDGSTS handling) to detect these opcodes and set insn_completed, allowing warp_inst_complete() to fire. Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com>
* simple_dram_cycle for loop * H200 DRAM cycle multiplier * Potential fix for pull request finding Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com> * Potential fix for pull request finding Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com> * Potential fix for pull request finding Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com> * formatting; H100 config mod * more opt * formatting * reverting gpgpu_n_mem * Simplify simple_dram_model_cycle arbitration loop Flatten nested loop structure into single loop with continue on DRAM full. Use processed counter to track requests issued per cycle. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> --------- Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com> Co-authored-by: JRPan <25518778+JRPan@users.noreply.github.com> Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com>
* Add NANOSLEEP and TRYWAIT retry support for spinlock simulation Implements non-blocking TRYWAIT semantics with configurable retry cycles (-gpgpu_trywait_max_retries, -gpgpu_trywait_retry_cycles). Adds NANOSLEEP_OP support to stall warps for specified durations. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Set gpgpu_chiplet_interleave to 0 for SM90 configurations * Refactor TRYWAIT to check in scheduler and only retry in replay regions Move TRYWAIT mbarrier check from issue_warp() to check_trywait_ready() called in the scheduler before issuing. This avoids ibuffer complications when the instruction needs to retry. - Add is_in_replay() virtual method to shd_warp_t (overridden in trace) - Only perform TRYWAIT retries when inside a replay region - Disable instruction prefetch during replay to avoid skipping REPLAY_END - Add SPECIALIZED_UNIT_5_OP to tensor core instruction counter - Change m_num_sim_winsn to uint64_t and expose as perf counter - XOR chiplet destination for interleaved mapping Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Track inflight TMA loads to block kernel from exiting * Format code * Add TMA warp tracking and mbarrier phase initialization - Add m_is_tma_warp flag to shd_warp_t for tracking warps with TMA instructions - Add first_trywait_checked flag and inc_phase() to mbarrier_t - Initialize mbarrier phase to 1 for TMA warps on first trywait check (nvjet/cuBLAS kernels only) - Modify retry logic: non-TMA warps retry indefinitely - Remove verbose debug print Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Add timestamped filenames for perf_counter output Generate timestamped perf_counter_YYYY-MM-DD_HH-MM-SS.csv.gz files and create a perf_counter.csv.gz symlink pointing to the latest file. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Simplify TRYWAIT retry condition Remove dead TMA warp check - all warps retry indefinitely for now, treating TRYWAIT as blocking in replay regions. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Optimize LRC coalescing stats to O(1) - Disable max_coalescing_count() which was iterating entire queue - Optimize avg_coalescing_count() to O(1) by tracking running total: - Increment on insert - Decrement by entry size on remove - Average = total / queue_size Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Reduce default max retry attempts for TRYWAIT from 100 to 5 * mark unused barriers and skip on TRYWAIT * Add STAS instruction support for Hopper - Add STAS_OP to op_type enum - Enable STAS in generate_mem_accesses() with address masking (lower 32 bits = store addr, upper 32 bits = mbarrier addr) - Add STAS_OP to ldst_unit::can_issue() and scheduler - Handle mbarrier_complete_tx when STAS store completes - Fix TRYWAIT phase check for nvjet kernels Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Ignore bar.sync in nvjet kernels. --------- Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com> Co-authored-by: Weili An <an107@purdue.edu>
…l-sim#23) * Add per-MC DRAM stats and fix L2_dram_queue_full counter location - Add dram_reads_per_mc and dram_writes_per_mc statistics counters to track DRAM traffic per memory controller - Move L2_dram_queue_full counter to l2_cache::cycle() where the DRAM queue fullness is actually checked (memport->full()) - Refactor cache_cycle() to remove redundant L2_dram_queue->full() check - Enable debug symbols (-g) and frame pointers in all builds for profiling - Enable chiplet interleave in H200 config Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Add CHIPLET_ACC access types for inter-chiplet traffic tracking Introduces CHIPLET_ACC_R and CHIPLET_ACC_W to separate inter-chiplet traffic from local GLOBAL_ACC traffic in statistics. Adds set_type() method to mem_access_t for reclassification when requests enter the chiplet interconnect queues. Also adjusts H100/H200 config clock domains and DRAM queue sizes. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Add mbarrier try_wait and nanosleep wait cycle tracking Add per-warp cycle recording for mbarrier TRYWAIT retries to analyze sync overhead and waiting patterns. Output as CSV format (MBARRIER_TRYWAIT,sid,warp_id,cycles...) in update_stats() for easy Python parsing. Also add per-core nanosleep_wait_cycles counter. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Add support for CHIPLET_ACC_R access type in atomic read assertions * Handle missing mbarrier in trywait and skip barrier for inactive warps Add lookup_clustermbar_allow_nonexist() to gracefully handle mbarrier lookup when the mbarrier hasn't been created yet. In trywait, retry with nanosleep if mbarrier is not found. Also skip BAR.SYNC handling when the warp has no active threads. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * disable STAS * Comment out debug symbols and frame pointers compilation options --------- Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com>
In issue_warp(), the barrier active_count check used next_inst->active_count(), but next_inst is the static instruction which has no active mask set in PTX mode. This caused all barriers to be silently skipped, breaking __syncthreads() synchronization (e.g. rodinia streamcluster). Use active_mask (from SIMT stack/trace) instead, which is the correct dynamic mask for both PTX and SASS modes. Co-authored-by: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
* Add PM workflow stubs that delegate to accel-sim-framework-private Three thin wrappers that reuse the workflows in purdue-aalp/accel-sim-framework-private@ai-automation instead of duplicating their bodies: - pr-ai-review.yml — on `ready` label / workflow_dispatch - issue-triage.yml — on issue opened / workflow_dispatch - pr-review-nudge.yml — Mon/Thu cron / workflow_dispatch Requires: - Repo secrets AALP_APP_ID, AALP_APP_PEM_PATH (already set). - Labels ready, ai-reviewed, ai-concerns, awaiting-human-review, stale-review, ai-triage (already created). - purdue-aalp/accel-sim-framework-private Actions access set to allow calls from this repo. - aalp-self-hosted-runner app installed here with Contents:Read + Pull requests:Write + Members:Read. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> * Copy AI PR review + review nudge workflows from accel-sim-framework Duplicate the two PM workflows (pr-ai-review.yml, pr-review-nudge.yml) and the shared skill (.github/skills/pr-review.md) into this repo. Kept separate rather than reusable across repos — simpler, no cross-repo plumbing. Requires (already set up): - Repo secrets AALP_APP_ID, AALP_APP_PEM_PATH - Labels ready, ai-reviewed, ai-concerns, awaiting-human-review, stale-review - aalp-self-hosted-runner app has Pull requests: Write + Members: Read Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> --------- Co-authored-by: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
) When an mbarrier ARRIVES instruction is issued after all prior LDGSTS instructions in the warp have already drained from m_pending_ldgsts, the assert(size > 0) fired. Treat the empty case as "nothing pending" by recording UID 0 (never present in the map) so writeback immediately completes the mbarrier. Co-authored-by: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
The accel-sim wrapper no longer scans traces at init time to populate trace-driven warp/mbarrier metadata. Remove the consumers that were either dead or unreachable on the .tracez path: - kernel_info_t::is_mbarrier_addr_used / register_used_mbarrier_addr and m_used_mbarrier_addrs (only writer was the removed scan) - mbarrier_t::is_used / set_used / m_is_used and the mbarrier_init warning that depended on is_mbarrier_addr_used; the "skip waiting if mbarrier has no ARRIVE" optimization (already commented out at the trywait site) is removed too - mbarrier_t::is_first_trywait_checked / mark_first_trywait_checked / m_first_trywait_checked, only used by the nvjet trywait workaround - trace_shd_warp_t::is_tma_warp / set_is_tma_warp / m_is_tma_warp and the nvjet first-trywait phase-bump workaround that depended on it Note: the nvjet first-trywait workaround is now gone, so nvjet kernels will deadlock at the first mbarrier trywait until a proper fix lands. Co-authored-by: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
The L2-clock block incremented gpu_stall_icnt2mem on every cycle where
a memory sub-partition input was full, regardless of whether ICNT
actually had a packet waiting for that sub-partition. During phases
with sparse or zero ICNT->L2 traffic but other reasons keeping the
sub-partition full, the counter climbed monotonically — reflecting
"L2 sub-partition full cycles" rather than real ICNT->L2 push stalls.
Add a lightweight icnt_has_packet(deviceID) peek to the interconnect
wrapper, backed by both interconnect backends:
- InterconnectInterface::HasPacket (intersim2) mirrors Pop's subnet/VC
iteration without mutating _boundary_buffer or _round_robin_turn.
- LocalInterconnect::HasPacket / xbar_router::Has_Packet check the
out_buffers non-empty.
- Wired through icnt_wrapper.{h,cc} following the existing
icnt_has_buffer pattern.
Guard the stall increment at gpu-sim.cc:2325 so it only fires when
ICNT actually has a packet pending for the sub-partition.
Co-authored-by: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
) (accel-sim#34) The 64-bit mbarrier state's pending-thread-count field has bit 63 as the phase-advance trigger. When hardware traces capture this bit set, trace-driven mbarrier init must reproduce the immediate phase advance or the simulated mbarrier waits for arrivals that already happened. Plumb a per-thread init_as_one flag through syncs_operand and pass the whole operand into ldst_unit::mbarrier_init so it can call inc_phase() on the freshly-constructed mbarrier when the flag is set. Co-authored-by: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
* Replace scoreboard hash/tree sets with fixed-size bitsets Use std::bitset<512> for reg_table and longopregs instead of std::unordered_set and std::set. Register numbers are bounded, so bitsets provide O(1) operations, better cache locality, and lower per-warp memory overhead (~64 bytes vs ~100+ for empty hash sets). Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Use shared_ptr for warp_inst_t in mem_fetch and add pool allocator Replace embedded warp_inst_t copy in mem_fetch with std::shared_ptr<warp_inst_t> to avoid deep-copying per-thread vectors and access queues into every memory request. Add a shared_ptr constructor overload so shader.cc allocation loops can share one warp_inst_t across multiple mem_fetch objects from the same instruction. Add a free-list pool allocator (operator new/delete overrides) to mem_fetch to reduce heap fragmentation from frequent allocation/deallocation of memory request packets. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> * Revert scoreboard bitset optimization to fix PTX mode crash The bitset<512> replacement for unordered_set/set in the scoreboard crashes in PTX execution-driven mode where virtual register numbers can exceed 512. PTX allocates register numbers sequentially across all types (int, float, pred), so complex kernels easily surpass this limit. Symptoms in CI: - heartwall: SIGABRT in Scoreboard::reserveRegister (regnum >= 512) - streamcluster: FUNC_TEST_FAILED (silent hazard check skip for registers >= 512 in checkCollision caused incorrect simulation) Revert to original unordered_set<unsigned> / set<unsigned> which handle unbounded register numbers. Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com> * Add wall_clock_ms timestamp to perf counter CSV Add a monotonic wall-clock timestamp (milliseconds) column to each row of the perf counter CSV output. This enables plot-perf-counters.py to compute runtime simulation rate (inst/sec and cycle/sec) over time, identifying program phases that are expensive to simulate. Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com> --------- Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com> Co-authored-by: yechen3 <liu2550@purdue.edu>
…ccel-sim#35) When a TMA load returns from global memory, ldst_unit::writeback() completed the mbarrier transaction count by access.get_size(). The coalescer rounds each access up to a full 32B sector, so a load whose in-bounds data only partially fills a sector (e.g. 16 of 32 bytes) completed 32 bytes. Meanwhile the OOB completion at issue time already accounts for tile - in_bounds bytes, which includes the unused bytes of that same sector. The shared bytes are therefore counted twice, driving tx_count negative (it overshoots zero), so the barrier's phase never advances and the consumer warp's TRYWAIT spins until deadlock detection aborts the simulation. Complete by the exact in-bounds bytes the access covers (access.get_byte_mask().count()) instead of the sector-rounded size. This restores the invariant sum(in-bounds bytes) + oob == tile for every alignment; sector-aligned loads are unaffected. Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
… extraction) (accel-sim#36) 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>
Brings in the two commits h100-test was missing: 03c1fe4 Add additive L2 stat for NVIDIA-comparable write hits (accel-sim#141) 7377472 Fix ptxinfo parser for CUDA 11+ ptxas resource-usage output (accel-sim#138) Clean auto-merge; gpu-cache.cc and gpu-cache.h were touched on both sides but merged without conflict. h100-test was already current with gpgpu-sim/gpgpu-sim_distribution dev. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
The trigger was commented out 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 PR into accel-sim/gpgpu-sim_distribution would run no CI at all. This restores the same trigger set upstream `dev` already uses. Co-Authored-By: Claude Opus 5 <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
5 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.
The performance-model half of Accel-Sim 2.0, from
purdue-aalp/gpgpu-sim_distribution-public:h100-test, proposed for merge intodev. 66 commits, 58 files.Paper: https://arxiv.org/abs/2608.22602
Pairs with the Accel-Sim framework PR
This is one half of a two-repo change and should land together with:
That PR carries the tracer, the SASS frontend, and the tooling; this one carries the timing model they drive. Neither builds against the other repo's current
devon its own.Both branches are synced with their respective upstream
devas of this PR — including #141 (additive L2 write-hit stat) and #138 (ptxinfo parser fix), merged in rather than rebased so no shared history was rewritten.What lands here
Hopper execution model. TMA engine (bulk tensor movement, bulk/store groups, CGA shared-memory multicast); async WGMMA with warpgroup commit/wait semantics and MMA latency that scales with the N tile dimension;
mbarrierproducer/consumer synchronization including async-proxy fences, remote arrive, and dynamictry_wait/ spinloop evaluation at replay time; 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.
Configs. Tested
SM90_H100/H200configurations.Review notes
mbarrier) are the substantive part; the rest is config and plumbing.