diff --git a/.claude/rules/ci-stats.md b/.claude/rules/ci-stats.md index 091634c427..a3d4f69c30 100644 --- a/.claude/rules/ci-stats.md +++ b/.claude/rules/ci-stats.md @@ -41,6 +41,10 @@ Read `docs/ci-stats.md` before redesigning any of this. per repo at a time (`concurrency`). No A100. - `--array` (with `--per-kernel`) logs tasks as `_`; keep job ids strings and `squeue -r` in job_status. Without `--array` nothing changes. +- Every CI SASS run (main-sass, main-tracer, weekly) launches with + `SIM_PER_KERNEL` (lib/common.sh: `--per-kernel --array -M 10G`) and verifies + with `--per-kernel-dirs`. Not main-docker: inside the image the launcher is + procman, and `--array` needs sbatch. - A trace depot is replaced only by `publish_depot` of a finished run. - `CI_SANDBOX=` is the only path override; do not add per-path env knobs. - Every simulated app needs hardware stats: `NO-HW-REFERENCE` fails diff --git a/.github/ci-notes.md b/.github/ci-notes.md index 691ddf45b6..2ca4b95e32 100644 --- a/.github/ci-notes.md +++ b/.github/ci-notes.md @@ -13,14 +13,16 @@ on a remote H200 over VPN) then SASS-H200 (`H200-SASS`: rodinia, ubench, TMA, cutlass_tma_small); Tracer-V100 (on tgrogers-dgx) then SASS-V100 (`QV100-SASS`: rodinia, ubench); SASS-H200-ML (`H200-SASS` on the hosted - `H200_ml_traces`, never traced here; `--per-kernel --array`: one Slurm array - task per kernel, job ids `_`). Each SASS job simulates, + ML traces of the static H200 depot, never traced here). Each SASS job simulates, correlates, archives. - gpgpu-sim is cloned at build time (`clone_gpgpusim` in `lib/common.sh`: the same-named branch of the code's gpgpu-sim fork if it has one, else upstream `dev`, also when there is no such fork; any other clone or fetch failure fails the build): a simulator failure may come from that repo, not this one. -- Simulations run as slurm jobs; the monitor prints each job's status. +- Simulations run as slurm jobs; the monitor prints each job's status. Every + SASS run except SASS-Simulation-Docker is per kernel (`--per-kernel --array + -M 10G`): one Slurm array task per kernel, job ids `_`, run dirs + `...//per-kernel//`. `FUNC_TEST_FAILED` = the simulation ran but its output was wrong; `FUNC_TEST_FAILED, ASSERT` = the simulator hit an assertion; `COMPLETE_ERR_FILE_HAS_CONTENTS` = completed but wrote to stderr (a warning diff --git a/.github/scripts/lib/common.sh b/.github/scripts/lib/common.sh index 85e5a9d0cb..e6ce8b1ece 100755 --- a/.github/scripts/lib/common.sh +++ b/.github/scripts/lib/common.sh @@ -186,7 +186,11 @@ build_accelsim() { # $2 - Config # $3 - Trace path # $4 - Job name suffix (default: $$) -# $5.. - Extra run_simulations.py flags (e.g. --per-kernel) +# $5.. - Extra run_simulations.py flags (e.g. -M) +# Every CI run is per kernel: one Slurm array task per kernel (get_stats +# --per-kernel-dirs collects them), so a long kernel neither serialises the +# app's other kernels nor delays its own failure. +SIM_PER_KERNEL=(--per-kernel --array -M 10G) run_simulations() { local benchmarks="$1" local config="$2" diff --git a/.github/scripts/main-sass.sh b/.github/scripts/main-sass.sh index 26b49ef145..cb2c2011f2 100755 --- a/.github/scripts/main-sass.sh +++ b/.github/scripts/main-sass.sh @@ -108,11 +108,11 @@ stage_simulate() { for run in "${RUNS[@]}"; do config="${run%%:*}" trace=$(trace_dir "$config") - run_simulations "${run#*:}" "$config" "$trace" "$JOB_NAME" + run_simulations "${run#*:}" "$config" "$trace" "$JOB_NAME" "${SIM_PER_KERNEL[@]}" done monitor_simulations "$JOB_NAME" 12 300 "stats-per-app-sass.csv" -v - verify_simulations "$JOB_NAME" "${RUNS[@]}" + verify_simulations --per-kernel-dirs "$JOB_NAME" "${RUNS[@]}" } stage_correlate() { diff --git a/.github/scripts/main-tracer.sh b/.github/scripts/main-tracer.sh index 1bdf1d0fb7..b671216e8d 100755 --- a/.github/scripts/main-tracer.sh +++ b/.github/scripts/main-tracer.sh @@ -115,7 +115,8 @@ stage_test() { -B "$BENCHMARKS" \ -C "$CONFIG" \ -T ./hw_run/traces/ \ - -N "$JOB_NAME" + -N "$JOB_NAME" \ + "${SIM_PER_KERNEL[@]}" monitor_simulations "$JOB_NAME" 12 300 "rodinia-stats-per-app.csv" -v } diff --git a/.github/scripts/weekly.sh b/.github/scripts/weekly.sh index 9243f726ae..97bd726bd6 100755 --- a/.github/scripts/weekly.sh +++ b/.github/scripts/weekly.sh @@ -8,8 +8,8 @@ # trace-h200 VPN, trace on the remote H200, publish (one Slurm job) # build-tracer-v100, trace-v100, tracer and apps on the runner; traces and # hw-stats-v100, publish-v100 hardware stats on gpu01 (gpu-lock sruns); publish -# simulate|correlate|archive run: h200, v100, or h200-ml (the hosted -# ML traces of H200_ml_traces, never traced here) +# simulate|correlate|archive run: h200, v100, or h200-ml (the ML traces +# in the static H200 depot, never traced here) # # Requires GITHUB_REPOSITORY and BRANCH_NAME (set by weekly.yml); the trace-h200 # stage also VPN_HOST, VPN_CERT and REMOTE_HOST (org secrets). @@ -23,7 +23,6 @@ source "$SCRIPT_DIR/lib/common.sh" H200_SUITES="rodinia_2.0-ft,rodinia-3.1,GPU_Microbenchmark,GPU_Microbenchmark_TMA,cutlass_tma_small" V100_SUITES="rodinia_2.0-ft,rodinia-3.1,GPU_Microbenchmark" ML_SUITES="llama_inference,llama_inference_2048,resnet_inference,bert_inference,dlrmv2_inference,unet3d_inference,retinanet_inference,sdxl_inference,mixtral_inference,deepseek_r1_inference,whisper_inference,qwen25_inference,qwen25_inference_2048,llama_training" -ML_DEPOT="$DEPOT/H200_ml_traces" TRACE_DEVICE=7 export OPENCONNECT_PATH="/home/tgrogers-raid/a/common/openconnect/openconnect" @@ -51,20 +50,16 @@ depot_dir() { echo "$SHARED_ROOT/nightly-$1/${GITHUB_REPOSITORY//\//_}" } -# Sets CONFIG, SUITES, TRACES, HW, SIM_FLAGS, STATS_FLAGS, MONITOR_HOURS and -# GROUP for run $1 +# Sets CONFIG, SUITES, TRACES, HW, MONITOR_HOURS and GROUP for run $1 gpu_vars() { - SIM_FLAGS=() STATS_FLAGS=() MONITOR_HOURS=14 GROUP=ubench + MONITOR_HOURS=14 GROUP=ubench case "$1" in h200) CONFIG="H200-SASS"; SUITES="$H200_SUITES" ;; v100) CONFIG="QV100-SASS"; SUITES="$V100_SUITES" ;; h200-ml) - # ~9K kernels, one array task each (accelsim-layout.sh in the depot - # adds the traces/ and device-0/ layout these paths read) + # ~9K kernels, in the static CI depot (docs/ci-stats.md) CONFIG="H200-SASS" SUITES="$ML_SUITES" GROUP=ml MONITOR_HOURS=40 - # one kernel needs far less than its app's accel-sim-mem; 10G ran all of them - SIM_FLAGS=(--per-kernel --array -M 10G) STATS_FLAGS=(--per-kernel-dirs) - TRACES="$ML_DEPOT/traces/device-0/12.8/" HW="$ML_DEPOT/" + HW="$DEPOT/H200/hw_run/" TRACES="$DEPOT/H200/hw_run/traces/device-7/12.8/" return ;; *) @@ -165,10 +160,10 @@ stage_simulate() { source ./env-setup/12.8_env_setup.sh source ./gpu-simulator/setup_environment.sh local job_name="weekly-$run-$(ci_job_name)" - run_simulations "$SUITES" "$CONFIG" "$TRACES" "$job_name" "${SIM_FLAGS[@]}" + run_simulations "$SUITES" "$CONFIG" "$TRACES" "$job_name" "${SIM_PER_KERNEL[@]}" # Weekly: one summary line per poll; a job the node's OOM killer took is not a failure monitor_simulations "$job_name" "$MONITOR_HOURS" 300 "weekly-$run-stats-per-app.csv" --allow-oom - verify_simulations "${STATS_FLAGS[@]}" --allow-oom "$job_name" "$CONFIG:$SUITES" + verify_simulations --per-kernel-dirs --allow-oom "$job_name" "$CONFIG:$SUITES" } stage_correlate() { diff --git a/.github/workflows/weekly.yml b/.github/workflows/weekly.yml index 8dd73eed1b..ac6da48fef 100644 --- a/.github/workflows/weekly.yml +++ b/.github/workflows/weekly.yml @@ -156,7 +156,7 @@ jobs: CI_ANALYSIS_FILES: "ci-stats/*/correl.txt ci-stats/*/*.kernel.raw.csv util/job_launching/logfiles/failed_job_log_*.txt" run: $AI/ci-analysis.sh - # The hosted ML traces of H200_ml_traces (never traced here), one Slurm array + # The ML traces of the static H200 depot (never traced here), one Slurm array # task per kernel; independent of the tracers SASS-H200-ML: runs-on: tgrogers-raid diff --git a/CITATION.cff b/CITATION.cff index 07516b973c..679060a0c7 100644 --- a/CITATION.cff +++ b/CITATION.cff @@ -26,7 +26,7 @@ message: >- type: software url: "https://accel-sim.github.io/" repository-code: "https://github.com/accel-sim/accel-sim-framework" -version: "2.0" +version: "2.0.1" # CFF requires a top-level `authors` key, so Accel-Sim is credited here as a single # entity rather than an enumerated list -- the project has many contributors beyond any # one paper's author list. See the GitHub contributors page for individual credit. diff --git a/README.md b/README.md index bbcbeaf06b..dc1a23d797 100644 --- a/README.md +++ b/README.md @@ -396,10 +396,16 @@ To disable it entirely, pass `--spinlock_handling none`: The detection tool lives in [`./util/tracer_nvbit/others/spinlock_tool/`](./util/tracer_nvbit/others/spinlock_tool/). **Pre-traced applications.** -To fetch a repository of pre-collected traces: +Pre-collected traces are published as GitHub release assets, one tar per GPU / workload / argument set, split into parts under 2 GB with SHA-256 checksums in the release's `manifest.json`. +To list, download, verify and extract them under `./hw_run//` in the standard layout (`traces/device-0////traces/` plus the Nsight Compute stats in `device-0////`): ```bash -./get-accel-sim-traces.py +./get-accel-sim-traces.py -l # list what is available +./get-accel-sim-traces.py -a H200/inference_bert.sh/32_256 # one argument set +./get-accel-sim-traces.py -a H200/inference_sdxl.sh,H200/all # GPU[/workload[/args]], 'all' at any level ``` +Pass `-T ./hw_run//traces/device-0/` to `run_simulations.py` and `-H ./hw_run/` to `plot-correlation.py`. +The release's `README.md` lists every set and how to download one by hand. +Downloading from a private repository needs `GITHUB_TOKEN` or a `gh auth login` session. ### SASS Frontend and Simulation Engine diff --git a/docs/ci-stats.md b/docs/ci-stats.md index 4fb3fd5c5c..c0e409e58c 100644 --- a/docs/ci-stats.md +++ b/docs/ci-stats.md @@ -6,7 +6,7 @@ |---|---|---|---| | pull_request (smoke) | QV100/H200: GPU_Microbenchmark_CI, rodinia_2.0-ft; H200: GPU_Microbenchmark_TMA_CI, cutlass_tma_CI, hopper_features_CI | no | no | | push to dev (full) | QV100/H200: GPU_Microbenchmark, rodinia_2.0-ft; H200: cutlass_tma_small, GPU_Microbenchmark_TMA | yes | yes | -| weekly | H200 (this week's traces): rodinia, ubench, TMA, cutlass_tma_small; QV100 (this week's V100 traces): rodinia, ubench; H200 ML (hosted `H200_ml_traces`, 14 suites, 20 runs, ~9K kernels) | all | all | +| weekly | H200 (this week's traces): rodinia, ubench, TMA, cutlass_tma_small; QV100 (this week's V100 traces): rodinia, ubench; H200 ML (static H200 depot, 14 suites, 20 runs, ~9K kernels) | all | all | Each SASS job: simulate (`run_simulations`), monitor, `verify_simulations` (`get_stats --strict` per config into `ci-stats//run.csv`), @@ -68,17 +68,17 @@ CSVs can be diffed by hand when needed. `weekly.sh` stage per step. No A100, no static-depot regression runs. - `concurrency: weekly-`: the depots are per repo, not per run, so a second run's tracer would replace the traces the first run's sims read. -- SASS-H200-ML replays `accel-sim/H200_ml_traces` (owned by another user; - its layout is `//{processed_traces,ncu/full,ncu/cycles}`). - `accelsim-layout.sh` in that depot adds, without touching its files, - `traces/device-0/12.8///traces` (links to `processed_traces`, with - its own `kernelslist.g`, which the depot's layout script may filter; the - hardware rows of a left-out kernel go too, so both sides pair one-to-one) and - `device-0/12.8///` with - `ncu.csv.nsight` (-> ncu_stats.csv) and `ncu.csv.gpc__cycles_elapsed.0` - (ncu_cycles.0.csv converted to the long one-metric format). Not in it: - `llama_inference_32k` (no traces), FA3 (`benchmark_hopper_forward_varlen_prefill.py`: - its 18 dirs hold only legacy ncu backups, no traces). +- SASS-H200-ML replays the ML traces in the static depot + `accel-sim/H200/hw_run` (`traces/device-7/12.8///traces`, HW in + `device-7/12.8///{ncu.csv.nsight,ncu.csv.gpc__cycles_elapsed.0}`), + copied 2026-10-20 from `accel-sim/H200_ml_traces` (another user's depot: + `//{processed_traces,ncu/full,ncu/cycles}`) by that depot's + `accelsim-layout.sh` layer: the kernels in `H200/ML-DROPPED-KERNELS.md` are + left out of `kernelslist.g` and their HW rows dropped too, so both sides pair + one-to-one; `ncu.csv.gpc__cycles_elapsed.0` is ncu_cycles.0.csv in the long + one-metric format. Not in it: `llama_inference_32k` (no traces), FA3 + (`benchmark_hopper_forward_varlen_prefill.py`: no traces). Like every + static-depot trace, nothing updates them automatically. - ~9K per-kernel jobs: `run_simulations.py --per-kernel --array` submits each app's kernels as Slurm job arrays (<= MaxArraySize = 1001 tasks), one queue record while pending, instead of ~9K sbatch calls against the cluster-wide diff --git a/get-accel-sim-traces.py b/get-accel-sim-traces.py index e2350f8fb7..85046539d7 100755 --- a/get-accel-sim-traces.py +++ b/get-accel-sim-traces.py @@ -1,213 +1,249 @@ #!/usr/bin/env python3 +"""Download pre-collected Accel-Sim SASS traces from a GitHub release. -from optparse import OptionParser -import re -import os -import math -import sys -from subprocess import Popen, STDOUT - -VERSION = "1.1.0" -WEB_DIRECTORY = "ftp://ftp.ecn.purdue.edu/tgrogers/accel-sim/traces/" -this_directory = os.path.dirname(os.path.realpath(__file__)) + "/" -sys.path.insert(0, os.path.join(this_directory, "util", "job_launching")) -import common - -millnames = ["", " K", " M", " G", " T"] +Each trace set (one GPU / workload / argument set) is a tar split into parts +smaller than 2 GB. The release's manifest.json lists every part with its +SHA-256; this script downloads the selected parts, verifies them, and extracts +the tar under (default ./hw_run) in the standard Accel-Sim layout: + /traces/device-0////traces/ run_simulations.py -T + /device-0////ncu.csv.* plot-correlation.py -H +Private repositories need a token: set GITHUB_TOKEN, or log in with `gh`. +""" -def getNumRaw(n): +import argparse +import hashlib +import json +import os +import shutil +import subprocess +import sys +import tarfile +import urllib.error +import urllib.request + +DEFAULT_REPO = "accel-sim/accel-sim-framework" +# Traces ship as assets of the code release they were validated with. +DEFAULT_TAG = "v2.0.1" +API = "https://api.github.com" +CHUNK = 1 << 22 +EXTRACT_KW = {"filter": "data"} if hasattr(tarfile, "data_filter") else {} +this_directory = os.path.dirname(os.path.realpath(__file__)) + + +def human(n): + for unit in ("B", "KB", "MB", "GB", "TB"): + if abs(n) < 1000 or unit == "TB": + return f"{n:.1f} {unit}" if unit != "B" else f"{n} B" + n /= 1000.0 + + +def get_token(): + token = os.environ.get("GITHUB_TOKEN") or os.environ.get("GH_TOKEN") + if token: + return token + if shutil.which("gh"): + r = subprocess.run(["gh", "auth", "token"], capture_output=True, text=True) + if r.returncode == 0 and r.stdout.strip(): + return r.stdout.strip() + return None + + +class NoRedirect(urllib.request.HTTPRedirectHandler): + def redirect_request(self, req, fp, code, msg, headers, newurl): + return None + + +def open_url(url, token=None, accept=None): + """Open url; with a token, follow the asset redirect without forwarding Authorization.""" + headers = {"User-Agent": "accel-sim-get-traces"} + if accept: + headers["Accept"] = accept + if token: + headers["Authorization"] = f"Bearer {token}" + opener = urllib.request.build_opener(NoRedirect) try: - return float(n) - except ValueError: - count = 0 - for name in millnames: - if n[-1].strip() == name.strip(): - n = float(n[:-1].strip()) * 10 ** (3 * count) - break - count += 1 - return float(n) - - -def millify(n): - n = getNumRaw(n) - if math.isnan(n): - return "NaN" - if math.isinf(n): - return "inf" - millidx = max( - 0, - min( - len(millnames) - 1, int(math.floor(0 if n == 0 else math.log10(abs(n)) / 3)) - ), - ) - return "{:.2f}{}".format(n / 10 ** (3 * millidx), millnames[millidx]) - - -class Suite: - def __init__(self, name): - self.name = name - self.uncompressedSize = None - self.compressedSize = None - - -class Card: - def __init__(self, name): - self.name = name - self.suites = {} - - def getTotalCompressed(self): - total = 0.0 - for name, suite in self.suites.items(): - try: - total += suite.compressedSize - except TypeError: - sys.exit("Problem with compressed size in suite {0}".format(name)) - return total - - def getTotalUncompressed(self): - total = 0.0 - for name, suite in self.suites.items(): - total += suite.uncompressedSize - return total - - -def downloadTrace(cardName, suiteName): - webFile = os.path.join( - WEB_DIRECTORY, cardName, VERSION + ".latest", suiteName + ".tgz" - ) - print("\n\nDownloading {0}".format(webFile)) - wget = Popen(["wget " + webFile], stderr=STDOUT, shell=True) - wget.communicate() - if wget.returncode != 0: - sys.exit("wget {0} returned {1}".format(webFile, wget.returncode)) + return opener.open(urllib.request.Request(url, headers=headers)) + except urllib.error.HTTPError as e: + if e.code in (301, 302, 303, 307, 308): + loc = e.headers["Location"] + return urllib.request.urlopen(urllib.request.Request(loc, headers={"User-Agent": headers["User-Agent"]})) + raise -def main(): - parser = OptionParser() - parser.add_option( - "-a", - "--apps", - dest="apps", - default=None, - help="Pass the comma seperated input list instead of asking " - "from stdin. Example: -a tesla-v100/rodinia-3.1,tesla-v100/cudasdk", - ) - parser.add_option( - "-d", - "--download_dir", - dest="download_dir", - default="", - help="Directory to download the traces to.", - ) - (options, args) = parser.parse_args() - try: - hw_run_dir = common.dir_option_test( - options.download_dir, "hw_run", this_directory - ) - except common.PathMissing: - hw_run_dir = os.path.join(this_directory, "hw_run") - - if not os.path.exists(hw_run_dir): - os.makedirs(hw_run_dir) - os.chdir(hw_run_dir) - # Parse the trace summary - trace_summary = os.path.join(hw_run_dir, VERSION + ".trace.summary.txt") - try: - os.remove(trace_summary) - except OSError: +class Release: + def __init__(self, repo, tag, token): + self.repo, self.tag, self.token = repo, tag, token + try: + with open_url(f"{API}/repos/{repo}/releases/tags/{tag}", token, "application/vnd.github+json") as r: + rel = json.load(r) + except urllib.error.HTTPError as e: + hint = "" if token else " (private repositories need GITHUB_TOKEN or `gh auth login`)" + sys.exit(f"Cannot read release {tag} of {repo}: HTTP {e.code}{hint}") + self.assets = {a["name"]: a for a in rel["assets"]} + + def open_asset(self, name): + if name not in self.assets: + sys.exit(f"Asset {name} is not in release {self.tag} of {self.repo}") + a = self.assets[name] + if self.token: + return open_url(a["url"], self.token, "application/octet-stream") + return open_url(a["browser_download_url"]) + + +def sha256_file(path): + h = hashlib.sha256() + with open(path, "rb") as f: + for b in iter(lambda: f.read(CHUNK), b""): + h.update(b) + return h.hexdigest() + + +def download_part(release, part, dest): + """Download one part to dest, skipping it if a verified copy is already there.""" + if os.path.exists(dest) and os.path.getsize(dest) == part["size"] and sha256_file(dest) == part["sha256"]: + print(f" {part['name']}: already downloaded") + return + tmp = dest + ".tmp" + h, done = hashlib.sha256(), 0 + with release.open_asset(part["name"]) as r, open(tmp, "wb") as f: + for b in iter(lambda: r.read(CHUNK), b""): + f.write(b) + h.update(b) + done += len(b) + print(f"\r {part['name']}: {human(done)} / {human(part['size'])}", end="", flush=True) + print() + if done != part["size"] or h.hexdigest() != part["sha256"]: + os.remove(tmp) + sys.exit(f"{part['name']}: size/SHA-256 mismatch (got {done} bytes, expected {part['size']})") + os.replace(tmp, dest) + + +class Concat: + """Read a list of files as one stream, hashing everything read.""" + + def __init__(self, paths): + self.paths, self.f, self.h = list(paths), None, hashlib.sha256() + + def read(self, n=-1): + out = b"" + while n < 0 or len(out) < n: + if self.f is None: + if not self.paths: + break + self.f = open(self.paths.pop(0), "rb") + b = self.f.read(-1 if n < 0 else n - len(out)) + if not b: + self.f.close() + self.f = None + continue + out += b + self.h.update(out) + return out + + +def set_name(e): + return f"{e['gpu']}/{e['workload']}/{e['args']}" + + +def extract(entry, part_paths, out_dir): + """Extract one set; every member must be a file or directory inside its trace or hw directory.""" + allowed = [d for d in (entry["trace_dir"], entry["hw_dir"]) if d] + stream = Concat(part_paths) + with tarfile.open(fileobj=stream, mode="r|") as tar: + root = os.path.realpath(out_dir) + for m in tar: + target = os.path.realpath(os.path.join(out_dir, m.name)) + if not target.startswith(root + os.sep) or not (m.isfile() or m.isdir()): + sys.exit(f"Refusing unexpected tar member {m.name}") + if not any(m.name == d or m.name.startswith(d + "/") for d in allowed): + sys.exit(f"Tar member {m.name} is outside {allowed}") + tar.extract(m, out_dir, **EXTRACT_KW) + while stream.read(CHUNK): # hash trailing padding so the whole-tar checksum covers every byte pass + if stream.h.hexdigest() != entry["sha256"]: + sys.exit(f"{set_name(entry)}: whole-tar SHA-256 mismatch after extraction") + listed = [l for l in open(os.path.join(out_dir, entry["trace_dir"], "kernelslist.g")) if l.startswith("kernel-")] + if len(listed) != entry["kernels"]: + sys.exit(f"{set_name(entry)}: kernelslist.g has {len(listed)} kernels, manifest says {entry['kernels']}") + + +def select(manifest, selection): + """selection: comma list of GPU[/workload[/args]], each level may be 'all'.""" + chosen = [] + for item in selection.split(","): + want = [x for x in item.strip().split("/") if x] + if not want: + continue + hits = [e for e in manifest["traces"] + if all(w in ("all", v) for w, v in zip(want, (e["gpu"], e["workload"], e["args"])))] + if not hits: + raise ValueError(f"nothing matches '{item}'") + chosen += [e for e in hits if e not in chosen] + return chosen - Popen( - ["wget " + " " + WEB_DIRECTORY + VERSION + ".trace.summary.txt"], - stderr=STDOUT, - shell=True, - ).communicate() - - lineFormat = re.compile(r"(.*)\t(.*)/" + VERSION + ".latest/(.*)") - sizeDict = {} - for line in open(trace_summary): - lineMatch = lineFormat.match(line) - if lineMatch != None: - size = lineMatch.group(1) - cardName = lineMatch.group(2) - fileName = lineMatch.group(3) - suiteName = re.sub(r"(.*)\.tgz", r"\1", fileName) - if cardName not in sizeDict: - sizeDict[cardName] = Card(cardName) - card = sizeDict[cardName] - if suiteName not in card.suites: - newSuite = Suite(suiteName) - card.suites[suiteName] = Suite(suiteName) - - suite = sizeDict[cardName].suites[suiteName] - if suiteName != fileName: - suite.compressedSize = getNumRaw(size) - else: - suite.uncompressedSize = getNumRaw(size) - - # Infor the user what is available - ask them what they want to do - print("\n\nCurrently Available Traces:") - for cardName, card in sizeDict.items(): - print( - "GPU Name: {0}. All Apps Compressed (Download size): {1}, All Apps Uncompressed (Size on disk when used): {2}".format( - card.name, - millify(card.getTotalCompressed()), - millify(card.getTotalUncompressed()), - ) - ) - for name, suite in card.suites.items(): - print( - "\t{0}: Compressed = {1}, Uncompressed = {2}".format( - suite.name, - millify(suite.compressedSize), - millify(suite.uncompressedSize), - ) - ) - - selectionValid = False - while not selectionValid: - if options.apps == None: - selection = input( - "\n-------\nWhat do you want to download?" - "\n, (i.e. tesla-v100/rodinia-3.1,tesla-v100/cudasdk)" - "\n(Default=all/all) : " - ) - if selection == "" or selection == None: - selection = "all/all" - else: - selection = options.apps + +def main(): + p = argparse.ArgumentParser(description=__doc__, formatter_class=argparse.RawDescriptionHelpFormatter) + p.add_argument("-a", "--apps", help="Comma separated GPU[/workload[/args]] list, 'all' allowed at any level. " + "Example: -a H200/inference_bert.sh/32_256,H200/inference_sdxl.sh") + p.add_argument("-d", "--download_dir", default=os.path.join(this_directory, "hw_run"), + help="Extract into // (default: ./hw_run)") + p.add_argument("-l", "--list", action="store_true", help="List the available traces and exit") + p.add_argument("--repo", default=DEFAULT_REPO, help=f"GitHub repository (default: {DEFAULT_REPO})") + p.add_argument("--tag", default=DEFAULT_TAG, help=f"Release tag (default: {DEFAULT_TAG})") + p.add_argument("--keep_parts", action="store_true", help="Keep the downloaded tar parts after extraction") + o = p.parse_args() + + release = Release(o.repo, o.tag, get_token()) + with release.open_asset("manifest.json") as r: + manifest = json.load(r) + assert manifest.get("format") == 2, f"unsupported manifest format {manifest.get('format')}, expected 2" + + print(f"Available traces in {o.repo} release {o.tag}:") + for e in sorted(manifest["traces"], key=lambda e: (e["gpu"], e["workload"], e["args"])): + print(f" {e['gpu']}/{e['workload']}/{e['args']}: {e['kernels']} kernels, {human(e['size'])}") + total = sum(e["size"] for e in manifest["traces"]) + print(f" total: {len(manifest['traces'])} trace sets, {human(total)}") + if o.list: + return + + while True: + selection = o.apps or input("\nWhat do you want to download? GPU[/workload[/args]],... (default all): ") or "all" try: - for item in selection.split(","): - cardName = item.split(r"/")[0] - suiteName = item.split(r"/")[1] - if cardName == "all": - for cardName, card in sizeDict.items(): - if suiteName == "all": - for suiteName, suite in card.suites.items(): - downloadTrace(cardName, suiteName) - else: - downloadTrace(cardName, suiteName) - else: - card = sizeDict[cardName] - if suiteName == "all": - for suiteName, suite in card.suites.items(): - downloadTrace(cardName, suiteName) - else: - downloadTrace(cardName, suiteName) - - selectionValid = True - except Exception as e: - selectionValid = False - print("Invalid Input: {0}".format(e)) - if options.apps != None: + chosen = select(manifest, selection) + break + except ValueError as e: + print(f"Invalid input: {e}") + if o.apps: sys.exit(1) - print( - "\n\nDownload successful to {0}.\nFiles must be uncompressed with tar -xzvf to be usable by accel-sim".format( - hw_run_dir - ) - ) + out_dir = os.path.abspath(o.download_dir) + parts_dir = os.path.join(out_dir, ".trace-parts") + os.makedirs(parts_dir, exist_ok=True) + print(f"\nDownloading {len(chosen)} trace sets, {human(sum(e['size'] for e in chosen))}, into {out_dir}") + for e in chosen: + print(f"\n{e['gpu']}/{e['workload']}/{e['args']} ({len(e['parts'])} parts)") + paths = [] + for part in e["parts"]: + dest = os.path.join(parts_dir, part["name"]) + download_part(release, part, dest) + paths.append(dest) + extract(e, paths, out_dir) + if not o.keep_parts: + for path in paths: + os.remove(path) + print(f" extracted to {os.path.join(out_dir, e['trace_dir'])}") + if not o.keep_parts: + try: + os.rmdir(parts_dir) + except OSError: + pass + print("\nDone.") + for t, h in sorted({(e["trace_root"], e["hw_root"]) for e in chosen}): + gpu_dir = os.path.join(out_dir, t.split("/")[0]) + print(f" run_simulations.py -T {os.path.join(out_dir, t)}") + print(f" plot-correlation.py -H {gpu_dir}") if __name__ == "__main__": diff --git a/release.notes.md b/release.notes.md index cd097d0ec9..b649ab47b6 100644 --- a/release.notes.md +++ b/release.notes.md @@ -1,4 +1,58 @@ +# Release v2.0.1 + +_October 2026._ + +A maintenance release on top of 2.0.0: pre-collected traces move from the FTP server to +GitHub release assets, trace post-processing and parsing fixes, and a reworked CI. + +> **Requires the matching GPGPU-Sim:** `gpgpu-sim_distribution` tag `v2.0.1`. + +## Pre-collected traces on GitHub releases +- Traces are now published as assets of the code release they were validated with, starting + with this one, so code and traces stay consistent. Each GPU / workload / argument set is one + tar, split into parts under 2 GB, and `manifest.json` records the SHA-256 of every part and + of the whole tar. The release also carries a `README.md` with download instructions. +- The first set is the **H200 ML suite**: 20 argument sets of BERT, ResNet-50, DLRMv2, + 3D U-Net, RetinaNet, Whisper, SDXL, Llama 3.1 8B (inference and training), Qwen 2.5 7B, + DeepSeek R1 and Mixtral 8x7B, traced with NVBit 1.8 as `.tracez`. Each argument set comes + with its Nsight Compute stats, row-aligned with `kernelslist.g`. A few kernels that do not + complete in simulation are left out of both. +- `get-accel-sim-traces.py` is rewritten. It lists the available sets (`-l`), downloads a + selection (`-a H200/inference_bert.sh/32_256`, with `all` at any level), checks every + checksum, and extracts under `hw_run//` in the standard layout, ready for + `run_simulations.py -T hw_run//traces/device-0/` and `plot-correlation.py -H hw_run/`. The FTP download path + and the pre-2.0 trace listing are removed. + +## Tracer +- `post-traces-processing` decodes every lane's `mbarrier` address when dropping unused + `TRYWAIT`s (#564). It used to read only the first lane of `base_stride` / `base_delta` + lines, so a `TRYWAIT` on an `mbarrier` reached only through another lane could be dropped + and the kernel could deadlock (#561). + +## Trace parsing +- Opening a `.tracez` file checks every `fread`, so a truncated file fails with an error + instead of filling the thread-block index with garbage. This also removes the framework's + remaining build warnings (#565). + +## Job launching +- `slurm.sim` is back to the pre-2.0 template (#556). + +## CI +- PRs run a smoke tier, every `dev` push runs correlated QV100 and H200 suites, and a + single Weekly workflow replaces the Volta and Hopper weeklies, adding the H200 ML suite + (#566). +- Issue triage and PR review automation moved to `purdue-aalp/github-ci`, with Slurm + cleanup and fork-aware fixes (#559). +- CI also runs on pushes to the purdue-aalp forks (#569). +- Every SASS CI run launches per kernel, as Slurm arrays at 10 GB, so one long kernel no + longer serialises the rest of its application. The Weekly H200 ML job reads its traces + from the static H200 CI depot. + +## GPGPU-Sim v2.0.1 +- Fixes 47 build warnings and two statistics that read uninitialized memory (#147). +- CI picks its runner pool by repository visibility (#144); README badge cleanup (#146). + # Release v2.0.0 — Accel-Sim 2.0 (full Hopper support) _Aug 25, 2026._