nki(mul2): NKI (Trainium) implementation - #294
Open
bowencui123 wants to merge 7 commits into
Open
bowencui123 wants to merge 7 commits into
bowencui123 wants to merge 7 commits into
Conversation
Split out of the consolidated NKI branch cecilia/feature/nki-vector-add (nki-all-operators, PR #259) so each operator can be reviewed on its own. Supersedes PR #185 (older per-operator branch). Co-Authored-By: Cecilia123li <68335867+Cecilia123li@users.noreply.github.com> Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_012Q38kGmXvyoeM1qtCbheSL
This was referenced Aug 29, 2026
Tunables mirror the Triton search space (`BLOCK_SIZE`); defaults are the previous constants, so autotune=False is unchanged. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_012Q38kGmXvyoeM1qtCbheSL
Merges NKI backend timing into results/csv/mul2_default.csv, run against this branch's impl_nki.py on trn2.3xlarge with the LNC2 execution contract (NEURON_LOGICAL_NC_CONFIG=2, NEURON_RT_NUM_CORES=1, NEURON_CC_FLAGS="--target trn2 --lnc 2"). All cases pass correctness verification. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01AQseF7nyesBh8KZAp8g7Cm
…); rerun default+autotune benchmarks Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01ScXYNjrrKGgDUVNHxv7HJt
bowencui123
added a commit
that referenced
this pull request
Sep 18, 2026
…ches
The open pull requests were written against the old source layout. This
moves any file still under a pre-refactor source directory onto the
tilebench package and rewrites the imports inside it, so each branch is a
merge plus one command:
git merge -c merge.directoryRenames=false origin/main
python scripts/migrate_pr_paths.py
Disabling directory-rename detection is load-bearing. With it on, git
guesses a destination for a file added inside a renamed directory and picks
the wrong one for the NKI operator PRs, offering
tilebench/benchmarks/llm_generated/<op>/<model>/high/final/impl_nki.py
instead of tilebench/benchmarks/operators/<op>/impl_nki.py.
Verified on #294 (single operator), #303 (cross-operator import) and #259
(45 operators at once). Temporary: delete it once those branches land.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA
Jokeren
pushed a commit
that referenced
this pull request
Sep 20, 2026
* chore(llm_codegen): move task descriptions next to the pipeline; drop benchmarks/problems benchmarks/problems/current/<op>_current.md is a prompt input of the LLM codegen pipeline, not a benchmark asset. Move the 45 descriptions to tools/llm_codegen/problems/ and repoint prompt_builder and generate_descriptions. The 16 unused files under problems/original are removed. Rebuilt iter-0 prompts keep a description section identical to the archived campaign prompts. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA (cherry picked from commit bbd8b61) * chore(results): stop tracking raw logs; add archive_logs.sh results/logs/ was force-added despite being git-ignored. Untrack the 272 raw timing/autotune JSON files; the summary CSVs stay tracked. scripts/archive_logs.sh snapshots the local results/logs/ onto archive/raw-logs-2026-09-18, building each commit as <origin/main tree> + results/logs/ in a throwaway index so that branch always equals main plus the raw logs. Nothing is checked out. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA (cherry picked from commit fec244e) * chore(peak_performance): move check_peak_specs.py beside the peak files Drop the .gitkeep placeholder (the directory holds B200.json and Trainium2.json). check_peak_specs.py now reads the measured peak bandwidth and per-dtype peak TFLOPS from B200.json instead of hardcoding the 8000 GB/s datasheet value. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA (cherry picked from commit 4b5f139) * data(llm_generated): restore the full two-model LLM campaign PR #98 cherry-picked the llm_codegen code without its data, leaving main with two pilot operators from #95. Restore the campaign from exp/llm_and_analysis: 44 operators x {gpt-5.5, claude-opus-4-7}/high, each with iter_N/ prompts, responses, kernels, feedback and token usage plus final/ and run_summary.json. Token counts match the archived RQ4 trajectory data iteration by iteration. Removed: the pipeline-development runs kept from #95 (flash_attention high_CHEAT / high_no_skill) and the vector_add directory. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA (cherry picked from commit 30ff150) * data(llm_generated): restore the vector_add campaign runs Bring back vector_add/{gpt-5.5,claude-opus-4-7}/high so the archive covers all 45 operators of the paper's LLM track. benchmarks/llm_generated is now identical to the tree on exp/llm_and_analysis. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA (cherry picked from commit d33b2fc) * feat(llm_codegen): ship the Triton and cuTile API references under skills/ prompt_builder read the API references from .claude/skills/, which is not tracked, so a fresh clone built prompts without them. Add skills/{triton,cutile}-guide/SKILL.md and read them from there. The files are the versions the paper's campaign used (the cuTile guide predates the later ct.static_iter section). Rebuilt iter-0 prompts now reproduce the API-reference and description blocks of all 90 archived prompts byte for byte. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA (cherry picked from commit a00dfa0) * feat(results): archive raw logs automatically after every run_bench run_bench.py is the only writer of results/logs/ (run_bench_all.py writes to results/runs/), so it now calls scripts/archive_logs.sh once the timing and autotune JSON are written. The snapshot is a local commit on the archive branch; nothing is checked out or pushed, a failure only prints a warning, and --no-archive skips it. archive_logs.sh now accumulates: the tree is origin/main + the logs already archived (remote and local tips) + results/logs/ from disk, with the disk winning. A machine holding only some of the logs therefore adds to the archive instead of replacing it. Offline runs fall back to the last fetched refs, and the ref update is compare-and-swap. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA (cherry picked from commit 3ebae04) * tools: add scripts/plot_sweep_max.py Builds the cross-operator sweep-max latency figure from results/csv and the NCU catalogue. Split out of the README work so the layout refactor, which rewrites this script's path handling, does not depend on a docs commit. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA * refactor: consolidate TileBench sources under the tilebench package Move the framework sources out of the repository root into one importable package, so imports are unambiguous and package resources resolve from the package location instead of the current working directory. core/ -> tilebench/core/ data/ -> tilebench/data/ benchmarks/ -> tilebench/benchmarks/ tools/llm_codegen/ -> tilebench/llm_codegen/ tilebench_run/ -> tilebench/profiling/ (source, catalogue) tilebench_run/ncu/ -> outputs/ncu/ (generated reports) New tilebench/paths.py resolves every package resource from __file__: operator configs, peak-performance JSON, LLM prompt inputs and the NCU catalogue. Callers that hardcoded relative paths (engine.py's benchmarks/operators/<op>/config.yaml), recomputed the repo root by parent count, or embedded this machine's absolute checkout path (ten profiling scripts) now go through it. scripts/ stay the user entry points and put the repository root on sys.path themselves, so they no longer need PYTHONPATH=. and work from any directory. Peak performance is split by kind: the framework-consumed device specs stay tracked at tilebench/data/peak_performance/<GPU>.json with unchanged contents, while the detailed measure_peak.py sweep becomes a generated output under the git-ignored outputs/peak_performance/ and is no longer tracked under results/. tilebench/profiling/hf_upload.py gains a __main__ guard: it uploads at import time, which was harmless for a loose script but not for a module of an importable package. Behaviour is unchanged. Verified: 45 operators discovered, all 45 config.yaml parse to identical hashes, all 180 impl_* modules import, the 45 rebuilt LLM prompts are byte-identical to the pre-move build, peak-performance and NCU catalogue loads match, and 6429 of 6507 moved files are byte-identical (the 78 modified carry import/path edits only). Committed with --no-verify so the whitespace hooks cannot rewrite 65 pre-existing files during a pure move. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA (cherry picked from commit 07c8540) * tools: add scripts/migrate_pr_paths.py for rebasing pre-refactor branches The open pull requests were written against the old source layout. This moves any file still under a pre-refactor source directory onto the tilebench package and rewrites the imports inside it, so each branch is a merge plus one command: git merge -c merge.directoryRenames=false origin/main python scripts/migrate_pr_paths.py Disabling directory-rename detection is load-bearing. With it on, git guesses a destination for a file added inside a renamed directory and picks the wrong one for the NKI operator PRs, offering tilebench/benchmarks/llm_generated/<op>/<model>/high/final/impl_nki.py instead of tilebench/benchmarks/operators/<op>/impl_nki.py. Verified on #294 (single operator), #303 (cross-operator import) and #259 (45 operators at once). Temporary: delete it once those branches land. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA (cherry picked from commit 0fa58d1) * chore: keep only CSV benchmark results tracked results/ now tracks only the summary CSVs. The per-operator figures, the aggregate tables and runtime_summary.md are regenerated by visualize.py and aggregate_results.py, so they leave version control; outputs/ was already declared the generated-artifact tree by the layout refactor but 49 NCU writeups and probe results were still tracked inside it. Nothing is deleted from the working tree: the 12 GB of .ncu-rep reports, the NCU writeups, the aggregates and the figures all stay on disk, now ignored. .gitignore previously ignored results/csv/ — exactly the one thing that should be version controlled — while the generated siblings were tracked by force-add. It now ignores results/* and re-includes results/csv/, verified with git check-ignore on existing, new and nested CSV paths. probe_kernel_count.py and ncu_writeup.py gained the mkdir their siblings already had: with outputs/ untracked, a fresh clone has no outputs/ncu/ to write into, and probe_kernel_count.py would otherwise fail after completing the whole probe. CSV contents are untouched: the 90 tracked files and their checksums are identical before and after. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA (cherry picked from commit b645387) * refactor: move kernel-count metadata into the profiling package kernel_counts.json records how many kernels one impl.run() launches per (op, dtype, backend) and under which names. The NCU harness validates every capture against it, so it is profiling input rather than a profiling report and belongs beside ncu_catalogue.json, not in the git-ignored outputs/ tree it shared with the 12 GB of .ncu-rep files. outputs/ncu/kernel_counts.json -> tilebench/profiling/kernel_counts.json The file is copied byte for byte (sha256 cfe0687e48665968e09e62ac3def2e86ae3bed14aa4a12a46175dcce79b666a6, 118850 bytes, 224 entries covering all 45 operators); no probe was re-run and no value was touched. tilebench.paths gains KERNEL_COUNTS, so every reader resolves it from the package instead of the working directory. Missing metadata is no longer silent. Both drivers used to skip the file when absent and then assume one kernel per launch, which disables wrong-kernel detection without saying so. ncu_kernel_select now owns the loader: an absent file raises MissingKernelCountsError, which the two CLIs turn into a clear message naming probe_kernel_count.py, while an individual unprobed pair still falls back to 1 and warns on stderr, since a single pair that errored during probing should not abort a whole sweep. probe_kernel_count.py writes the canonical path. A full run rewrites it as before; an ONLY_OP run now merges that operator's rows into the existing file, which previously truncated it to the probed operator alone. The probing algorithm, operator set, launch-count detection and JSON schema are unchanged. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA (cherry picked from commit c153413) * chore: stop tracking LLM-generated artifacts; add archive and fetch tooling tilebench/benchmarks/llm_generated/ stays the LLM pipeline's default output directory, but its contents are experiment artifacts rather than source, so main now git-ignores it. LLM_GENERATED_ROOT and generate.py are unchanged: same runtime path, different storage policy. The 6140 files are removed from the index only; nothing is deleted from the working tree. Backup: scripts/archive_artifacts.sh --logs | --llm | --all [--push] generalises archive_logs.sh. It still builds the commit in a throwaway index without checking anything out, and the archive is cumulative: every archive-owned path (results/logs/, tilebench/benchmarks/llm_generated/ and the pre-refactor benchmarks/llm_generated/ snapshot) is carried forward from the archive tip, so a machine holding only part of the artifacts adds to the backup instead of replacing it. archive_logs.sh remains as a wrapper for --logs, so run_bench.py keeps archiving raw logs only and never commits LLM trajectories as a side effect of a benchmark. Files are staged with update-index rather than git add -f: exclude pathspecs on ignored directories misbehave on git 2.27, and interpreter caches must stay out. Distribution: artifacts/manifest.json declares downloadable artifacts (URL, SHA256, archive name, provided directories). scripts/fetch_artifacts.py downloads one, verifies the checksum, unpacks into a staging directory and moves it into place, refusing to replace a non-empty destination without --force and rejecting archives with traversal, links or members outside the declared directories. Standard library only; Google Drive share links are rewritten to direct downloads. scripts/package_artifacts.py builds the reproducible tarball and prints the checksum for the manifest. The llm-aacl2026 entry keeps explicit placeholders for URL and SHA256 until the upload exists, and the downloader reports that instead of guessing. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA (cherry picked from commit 5c74683ebfaeba81268a1bcb665a3a8233ab9d82) * chore(artifacts): publish the llm-aacl2026 download URL and checksum The paper's LLM artifact is uploaded, so the manifest placeholders become the file's Google Drive share link and the SHA256 of the packaged archive. A real download through fetch_artifacts.py verified the checksum and restored 6140 files identical to the source tree. The manifest test now accepts either a fully published entry or an explicit placeholder, never a partial one. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA (cherry picked from commit 042f6d8d0460cc31cc0724cccf0d75a84b797fda) * refactor(results): move the B200 summary CSVs to results/B200/csv Path-only migration ahead of hardware-scoped results. All 90 committed CSVs (45 operators x {default, autotune}) are B200 measurements, so they move from results/csv/ to results/B200/csv/. Pure rename: no file is regenerated, reformatted or newline-normalised (56 of them use CRLF). Committed with --no-verify so the whitespace hooks cannot touch the data. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA * feat(results): scope results by hardware, results/<gpu>/... TileBench will hold results for several GPUs (B200 today; GH200 and an AMD part to follow), and a global results/{csv,logs,figures,aggregate,runs} lets the same operator overwrite itself across machines. Everything is now results/<gpu>/{csv,logs,figures,aggregate,runs}; only csv/ is tracked. - tilebench/paths.py: results_root / results_{csv,logs,figures,aggregate, runs}_dir and hardware_label(), which accepts any single safe path component. There is no device list: a new GPU needs a label, not a patch. - run_bench.py, run_bench_all.py, visualize.py, plot_sweep_max.py, aggregate_results.py, ncu_catalogue.py take --gpu. It is required and has no default, so a run on another machine cannot land in B200's namespace. run_bench.py prints the namespace next to the detected device and warns when the label does not appear in the device name. - A GPU CSV holds measurements from that GPU only. The TileLang merge is unchanged but confined to results/<gpu>/csv/. NKI (AWS Trainium) can no longer be merged into, or run together with, GPU backends: it is selected alone, without --gpu, and writes to outputs/nki/. Its profiler artifacts move there from results/logs/ too. The default backend set, and 'all', is triton, cutile, tilelang. - plot_sweep_max.py writes to results/<gpu>/figures/ by default so plotting another GPU cannot overwrite the README figure; pass --output for that. - .gitignore tracks results/<hardware>/csv/**/*.csv for any label and keeps logs, figures, aggregates (CSV too) and runs ignored. - archive_artifacts.sh --logs takes --gpu and snapshots results/<gpu>/logs/ only. Logs of every other GPU on the archive branch, and the legacy results/logs/ and benchmarks/llm_generated/ snapshots, are carried forward. run_bench.py archives the logs of its own --gpu. - Tests: path validation, per-GPU defaults for every script, no collision between GPUs, TileLang merge isolation, NKI isolation, .gitignore rules, archive behaviour, and the 90 B200 CSVs checked against SHA256 recorded before the move. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA * fix(results): keep NKI results inside the campaign namespace The previous commit isolated NKI under outputs/nki/. That is reverted: NKI keeps its original result organisation, moved into the hardware-scoped layout. `run_bench.py --gpu B200 --tile-language nki` is allowed again; --gpu names the result namespace of the campaign and does not claim that NKI ran on that GPU. NKI still runs through its own path on AWS Trainium, and engine and kernel execution are unchanged. - The NKI merge into results/<gpu>/csv/<op>_{default,autotune}.csv is restored as it is on main: torch_nki_ms, nki_ms and speedup_nki are appended as measured, with no B200 drift scaling, and speedup_nki = torch_nki_ms / nki_ms, never torch_ms / nki_ms. - NKI timing and autotune JSON go to results/<gpu>/logs/ like every other log. The Neuron profiling artifacts move from results/logs/ to results/<gpu>/logs/nki_profiles/ and nki_neff_manifest.jsonl: the engine takes logs_dir and hands both paths to the orchestrator, which no longer has a default location. - --gpu is required again for every run. NKI still only runs when named: the default selection and 'all' stay triton, cutile, tilelang. run_bench_all.py goes back to the engine's own backend selection. - archive_artifacts.sh --logs --gpu <gpu> snapshots the whole logs/ tree, NKI profiles included. There is no separate NKI archive path. - Tests: NKI merges into the B200 CSV, its columns survive a later TileLang merge, frozen GPU columns and the GH200 CSV stay untouched, speedup_nki uses torch_nki_ms, NKI logs and profiles land under results/<gpu>/logs/ and are archived with it. The B200 guard now accepts the NKI extension: the manifest records the pre-move SHA256 and a digest of the frozen columns, so a file is either byte-identical or extended with its GPU columns intact. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA * refactor(profiling): move the B200 NCU metadata to metadata/B200/ Path-only move ahead of hardware-scoped profiling metadata. The catalogue (sweep-max cases and autotune winners) and the probed kernel counts were recorded on B200, so they move from tilebench/profiling/ to tilebench/profiling/metadata/B200/. No content change. Committed with --no-verify: neither file ends in a newline, and the end-of-file hook would otherwise rewrite them. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA * feat: multi-hardware consistency for logs, NCU metadata and the archive Four places still assumed a single GPU or a single run per operator. Raw JSON. results/<gpu>/logs/ held one <op>_results.json and one <op>_autotune.json, so a default run, an autotune run, a TileLang-only run and an NKI run overwrote each other. The default names now carry the mode and the backend selection: <op>_<mode>_<backends>.json, built only by tilebench.paths.timing_log_path / autotune_log_path. tilebench/backends.py owns the backend names, their canonical order and the --tile-language parser, so `cutile,triton` and `triton,cutile` are one tag. Explicit --output / --autotune-log still win. With one file per selection the autotune-log merge is gone: it existed to stop a TileLang run clobbering the Triton/cuTile winners, and would now refuse a rerun with other cases. visualize.py names the run it reads with --mode and --tile-language, and ncu_catalogue.py reads exactly <op>_autotune_<backends>.json (default triton,cutile; both are required). Nothing is chosen by glob or mtime. NCU metadata. The catalogue and the kernel counts were one global pair, although sweep-max cases, autotune winners, launch counts and kernel names differ between GPUs. They are now tilebench/profiling/metadata/<gpu>/{ncu_catalogue,kernel_counts}.json (the B200 files moved in the previous commit), reached through ncu_catalogue_path / kernel_counts_path. NCU_CATALOGUE and KERNEL_COUNTS are removed. ncu_catalogue, probe_kernel_count, ncu_one, ncu_driver, ncu_writeup, hf_upload and plot_sweep_max take --gpu, required and without a default. A GPU without metadata is an error that names the command to produce it; another GPU's files are never used. Reports go to outputs/ncu/<gpu>/, and hf_upload uploads under ncu_report_main/<gpu>/ so it cannot overwrite the released B200 reports. package-data covers the nested directory, checked by a non-editable install. Archive. main still tracks the legacy results/logs/, so reading main's tree would bring it back on every run. It is now dropped from the base tree, and results/logs/ on an archive tip is re-rooted at results/B200/logs/ byte for byte. A migrated file that meets a different file at its destination, already archived or in the working directory, aborts the run with nothing committed. The logs of other GPUs and both LLM paths are carried forward as before. The summary CSVs come from main's tree and are never copied. Portability. The profiling shell scripts located the repository through a hard-coded /projects/... path, a user-specific conda path and QOS, and PYTHONPATH=. They resolve it from their own location (batch2.sbatch from the submit directory) and work from any directory. run_batch.sh no longer renames or copies results: run_bench.py owns the names, and the <op>_summary.csv it copied has not existed for a while. run_bench.py starts the archive script inside the repository whatever the caller's directory. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA * Clean profiling package layout tilebench/ is the installable Python source. tilebench/profiling/ still held cluster launch scripts and measured data, which are neither. - Remove batch1_launch.sh, batch2.sbatch, rerun_timeout.sh, run_batch.sh and supplemental_launch.sh. They are operational scripts of one campaign on one cluster; no source file, script or test of the framework depends on them. They are kept on the archive branch, not moved to scripts/. - Move the NCU metadata out of the package. ncu_catalogue.json (autotune winners) and kernel_counts.json (probed launch counts and kernel names) are measured on one GPU and written by ncu_catalogue.py / probe_kernel_count.py: generated experiment data, not a package resource. No benchmark code path reads them; only the NCU tools do. They now live under the Git-ignored outputs/profiling/<gpu>/, next to the NCU reports, still one directory per GPU and still with no fallback to another GPU. The committed B200 files leave the tree and are kept on the archive branch. - plot_sweep_max.py no longer needs that metadata: the sweep-max case comes from the operator's config.yaml through the same rule the catalogue uses (ncu_catalogue.sweep_max_cases), so the README figure is reproducible from the configs and the committed CSVs alone. - pyproject.toml: drop the tilebench.profiling package-data entry. Wheel and sdist now carry only .py files for that package. - Tests: drop the two tests that only exercised the removed shell scripts, and the SHA256 fingerprints of the B200 CSVs and of the B200 metadata (tests/data/ is gone). They proved one migration, not a software contract. The CSV schema test, the namespace and isolation tests and the NKI merge tests stay. A packaging test now asserts that a regular install ships no shell script and no data under tilebench/profiling/. - archive_artifacts.sh: carry forward outputs/profiling/ and any launch script kept on the archive tip, and keep the tip's own tree as the base when it already contains origin/main, so code merged into the archive from a branch ahead of main is not rolled back by the next archive run. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA * Separate profiling tools from the profiling library tilebench/profiling/ still mixed three kinds of Python file. "All .py" is not the same as "package source", so each of the 17 files was classified by what it does and who imports it. Library, stays in the package (3): __init__.py ncu_kernel_select.py kernel selection, capture validation and metadata loading, imported by four tools ncu_catalogue.py sweep-max cases and catalogue entries; imported by plot_sweep_max.py. Its argparse main and sys.path hack are gone; write_catalogue() is the entry point. Command-line tools, moved to scripts/ (8): scripts/aggregate_results.py aggregates result CSVs, so it sits with visualize.py, not with NCU scripts/profiling/ncu_catalogue.py thin CLI over the library scripts/profiling/{ncu_driver,ncu_one,ncu_writeup,probe_kernel_count, hf_upload}.py had no importer outside the tests scripts/profiling/ncu_generic_harness.py never imported: it is the process NCU profiles, started by path They bootstrap sys.path like the other scripts and run from any directory without PYTHONPATH. The drivers locate the harness next to themselves instead of inside the package. The report path rule both drivers duplicated is now tilebench.paths.ncu_report_path. Campaign-specific, removed from the public tree (7), kept on the archive: probe_hung_cfgs.py and six ncu_harness_*.py. Hard-coded operators, cases and configs, code that runs at import, the pre-profiler-range methodology the generic harness replaced, and no reference from any driver, test or document. NCU behaviour, metadata and output paths are unchanged: a real capture of mul2 on a B200, started from /tmp with the relocated tools, validated 1/1 launches. archive_artifacts.sh now carries forward whatever the archive tip keeps under tilebench/profiling/ that the base tree dropped, not only *.sh/*.sbatch. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA * Decouple run_bench.py from the archive workflow A benchmark run only produces files. By default run_bench.py also fetched origin/main and the 7,000-file archive branch, and committed a snapshot of the raw logs to a local branch of the user's clone; the README had to tell readers to pass --no-archive on every command. Backing results up is a maintainer task that targets one dated archive branch, not a runtime feature. - run_bench.py: remove _archive_logs(), the --no-archive flag and the call after the run, with the subprocess and REPO_ROOT imports only they used. There is no opt-in replacement and no compatibility no-op. - Remove scripts/archive_artifacts.sh, scripts/archive_logs.sh and tests/test_archive_scripts.py. They are kept on the archive branch, where archive_artifacts.sh remains the maintenance tool of that branch. - results/<hardware>/logs/ stays Git-ignored: a run writes its logs there and performs no Git operation. - Tests: drop the archive hook from the run_bench fixture and the test of the hook itself; remove --no-archive from every call. A new test asserts that a run starts no process and that --no-archive and --archive are rejected. Namespace isolation, NKI merge, path layout and raw JSON collision tests are unchanged. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA * tests: the no-side-effect test names no removed flag Keep the assertion that matters, that run_bench.py starts no process, and drop the check that two archive flags are rejected: argparse rejects any unknown flag, and the public tree should not name an option it never had. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA * refactor: move problems/ to tilebench/problems and rename llm_codegen to llm Path-only. tilebench/llm_codegen/problems/ becomes tilebench/problems/, a sibling of benchmarks/, core/ and data/, and tilebench/llm_codegen/ becomes tilebench/llm/. No file content changes; the next commit updates the references. Committed with --no-verify: the task descriptions are prompt inputs and must stay byte-identical, so the whitespace hooks must not touch them. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA * Point the code at tilebench/problems and tilebench/llm Follows the path-only move. tilebench.paths gains PROBLEMS_ROOT and renames LLM_CODEGEN_ROOT to LLM_ROOT; prompt_builder.py and generate_descriptions.py read the task descriptions through PROBLEMS_ROOT instead of a directory next to themselves. Module references become tilebench.llm, run as `python -m tilebench.llm.generate`. tilebench/problems/ holds only Markdown, so pyproject.toml declares it as data of the top-level package, and migrate_pr_paths.py maps main's benchmarks/problems/current/ and tools/llm_codegen/ to the new locations. The 45 rebuilt initial prompts are byte-identical to the ones built before the move. A wheel and an sdist carry the 45 task descriptions and the llm package; tests assert that every operator has a description under PROBLEMS_ROOT and that a regular install ships them. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA --------- Co-authored-by: Claude Fable 5.1 <noreply@anthropic.com>
Jokeren
pushed a commit
that referenced
this pull request
Sep 20, 2026
* chore(llm_codegen): move task descriptions next to the pipeline; drop benchmarks/problems benchmarks/problems/current/<op>_current.md is a prompt input of the LLM codegen pipeline, not a benchmark asset. Move the 45 descriptions to tools/llm_codegen/problems/ and repoint prompt_builder and generate_descriptions. The 16 unused files under problems/original are removed. Rebuilt iter-0 prompts keep a description section identical to the archived campaign prompts. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA (cherry picked from commit bbd8b61) * chore(results): stop tracking raw logs; add archive_logs.sh results/logs/ was force-added despite being git-ignored. Untrack the 272 raw timing/autotune JSON files; the summary CSVs stay tracked. scripts/archive_logs.sh snapshots the local results/logs/ onto archive/raw-logs-2026-09-18, building each commit as <origin/main tree> + results/logs/ in a throwaway index so that branch always equals main plus the raw logs. Nothing is checked out. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA (cherry picked from commit fec244e) * chore(peak_performance): move check_peak_specs.py beside the peak files Drop the .gitkeep placeholder (the directory holds B200.json and Trainium2.json). check_peak_specs.py now reads the measured peak bandwidth and per-dtype peak TFLOPS from B200.json instead of hardcoding the 8000 GB/s datasheet value. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA (cherry picked from commit 4b5f139) * data(llm_generated): restore the full two-model LLM campaign PR #98 cherry-picked the llm_codegen code without its data, leaving main with two pilot operators from #95. Restore the campaign from exp/llm_and_analysis: 44 operators x {gpt-5.5, claude-opus-4-7}/high, each with iter_N/ prompts, responses, kernels, feedback and token usage plus final/ and run_summary.json. Token counts match the archived RQ4 trajectory data iteration by iteration. Removed: the pipeline-development runs kept from #95 (flash_attention high_CHEAT / high_no_skill) and the vector_add directory. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA (cherry picked from commit 30ff150) * data(llm_generated): restore the vector_add campaign runs Bring back vector_add/{gpt-5.5,claude-opus-4-7}/high so the archive covers all 45 operators of the paper's LLM track. benchmarks/llm_generated is now identical to the tree on exp/llm_and_analysis. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA (cherry picked from commit d33b2fc) * feat(llm_codegen): ship the Triton and cuTile API references under skills/ prompt_builder read the API references from .claude/skills/, which is not tracked, so a fresh clone built prompts without them. Add skills/{triton,cutile}-guide/SKILL.md and read them from there. The files are the versions the paper's campaign used (the cuTile guide predates the later ct.static_iter section). Rebuilt iter-0 prompts now reproduce the API-reference and description blocks of all 90 archived prompts byte for byte. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA (cherry picked from commit a00dfa0) * feat(results): archive raw logs automatically after every run_bench run_bench.py is the only writer of results/logs/ (run_bench_all.py writes to results/runs/), so it now calls scripts/archive_logs.sh once the timing and autotune JSON are written. The snapshot is a local commit on the archive branch; nothing is checked out or pushed, a failure only prints a warning, and --no-archive skips it. archive_logs.sh now accumulates: the tree is origin/main + the logs already archived (remote and local tips) + results/logs/ from disk, with the disk winning. A machine holding only some of the logs therefore adds to the archive instead of replacing it. Offline runs fall back to the last fetched refs, and the ref update is compare-and-swap. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA (cherry picked from commit 3ebae04) * tools: add scripts/plot_sweep_max.py Builds the cross-operator sweep-max latency figure from results/csv and the NCU catalogue. Split out of the README work so the layout refactor, which rewrites this script's path handling, does not depend on a docs commit. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA * refactor: consolidate TileBench sources under the tilebench package Move the framework sources out of the repository root into one importable package, so imports are unambiguous and package resources resolve from the package location instead of the current working directory. core/ -> tilebench/core/ data/ -> tilebench/data/ benchmarks/ -> tilebench/benchmarks/ tools/llm_codegen/ -> tilebench/llm_codegen/ tilebench_run/ -> tilebench/profiling/ (source, catalogue) tilebench_run/ncu/ -> outputs/ncu/ (generated reports) New tilebench/paths.py resolves every package resource from __file__: operator configs, peak-performance JSON, LLM prompt inputs and the NCU catalogue. Callers that hardcoded relative paths (engine.py's benchmarks/operators/<op>/config.yaml), recomputed the repo root by parent count, or embedded this machine's absolute checkout path (ten profiling scripts) now go through it. scripts/ stay the user entry points and put the repository root on sys.path themselves, so they no longer need PYTHONPATH=. and work from any directory. Peak performance is split by kind: the framework-consumed device specs stay tracked at tilebench/data/peak_performance/<GPU>.json with unchanged contents, while the detailed measure_peak.py sweep becomes a generated output under the git-ignored outputs/peak_performance/ and is no longer tracked under results/. tilebench/profiling/hf_upload.py gains a __main__ guard: it uploads at import time, which was harmless for a loose script but not for a module of an importable package. Behaviour is unchanged. Verified: 45 operators discovered, all 45 config.yaml parse to identical hashes, all 180 impl_* modules import, the 45 rebuilt LLM prompts are byte-identical to the pre-move build, peak-performance and NCU catalogue loads match, and 6429 of 6507 moved files are byte-identical (the 78 modified carry import/path edits only). Committed with --no-verify so the whitespace hooks cannot rewrite 65 pre-existing files during a pure move. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA (cherry picked from commit 07c8540) * tools: add scripts/migrate_pr_paths.py for rebasing pre-refactor branches The open pull requests were written against the old source layout. This moves any file still under a pre-refactor source directory onto the tilebench package and rewrites the imports inside it, so each branch is a merge plus one command: git merge -c merge.directoryRenames=false origin/main python scripts/migrate_pr_paths.py Disabling directory-rename detection is load-bearing. With it on, git guesses a destination for a file added inside a renamed directory and picks the wrong one for the NKI operator PRs, offering tilebench/benchmarks/llm_generated/<op>/<model>/high/final/impl_nki.py instead of tilebench/benchmarks/operators/<op>/impl_nki.py. Verified on #294 (single operator), #303 (cross-operator import) and #259 (45 operators at once). Temporary: delete it once those branches land. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA (cherry picked from commit 0fa58d1) * chore: keep only CSV benchmark results tracked results/ now tracks only the summary CSVs. The per-operator figures, the aggregate tables and runtime_summary.md are regenerated by visualize.py and aggregate_results.py, so they leave version control; outputs/ was already declared the generated-artifact tree by the layout refactor but 49 NCU writeups and probe results were still tracked inside it. Nothing is deleted from the working tree: the 12 GB of .ncu-rep reports, the NCU writeups, the aggregates and the figures all stay on disk, now ignored. .gitignore previously ignored results/csv/ — exactly the one thing that should be version controlled — while the generated siblings were tracked by force-add. It now ignores results/* and re-includes results/csv/, verified with git check-ignore on existing, new and nested CSV paths. probe_kernel_count.py and ncu_writeup.py gained the mkdir their siblings already had: with outputs/ untracked, a fresh clone has no outputs/ncu/ to write into, and probe_kernel_count.py would otherwise fail after completing the whole probe. CSV contents are untouched: the 90 tracked files and their checksums are identical before and after. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA (cherry picked from commit b645387) * refactor: move kernel-count metadata into the profiling package kernel_counts.json records how many kernels one impl.run() launches per (op, dtype, backend) and under which names. The NCU harness validates every capture against it, so it is profiling input rather than a profiling report and belongs beside ncu_catalogue.json, not in the git-ignored outputs/ tree it shared with the 12 GB of .ncu-rep files. outputs/ncu/kernel_counts.json -> tilebench/profiling/kernel_counts.json The file is copied byte for byte (sha256 cfe0687e48665968e09e62ac3def2e86ae3bed14aa4a12a46175dcce79b666a6, 118850 bytes, 224 entries covering all 45 operators); no probe was re-run and no value was touched. tilebench.paths gains KERNEL_COUNTS, so every reader resolves it from the package instead of the working directory. Missing metadata is no longer silent. Both drivers used to skip the file when absent and then assume one kernel per launch, which disables wrong-kernel detection without saying so. ncu_kernel_select now owns the loader: an absent file raises MissingKernelCountsError, which the two CLIs turn into a clear message naming probe_kernel_count.py, while an individual unprobed pair still falls back to 1 and warns on stderr, since a single pair that errored during probing should not abort a whole sweep. probe_kernel_count.py writes the canonical path. A full run rewrites it as before; an ONLY_OP run now merges that operator's rows into the existing file, which previously truncated it to the probed operator alone. The probing algorithm, operator set, launch-count detection and JSON schema are unchanged. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA (cherry picked from commit c153413) * chore: stop tracking LLM-generated artifacts; add archive and fetch tooling tilebench/benchmarks/llm_generated/ stays the LLM pipeline's default output directory, but its contents are experiment artifacts rather than source, so main now git-ignores it. LLM_GENERATED_ROOT and generate.py are unchanged: same runtime path, different storage policy. The 6140 files are removed from the index only; nothing is deleted from the working tree. Backup: scripts/archive_artifacts.sh --logs | --llm | --all [--push] generalises archive_logs.sh. It still builds the commit in a throwaway index without checking anything out, and the archive is cumulative: every archive-owned path (results/logs/, tilebench/benchmarks/llm_generated/ and the pre-refactor benchmarks/llm_generated/ snapshot) is carried forward from the archive tip, so a machine holding only part of the artifacts adds to the backup instead of replacing it. archive_logs.sh remains as a wrapper for --logs, so run_bench.py keeps archiving raw logs only and never commits LLM trajectories as a side effect of a benchmark. Files are staged with update-index rather than git add -f: exclude pathspecs on ignored directories misbehave on git 2.27, and interpreter caches must stay out. Distribution: artifacts/manifest.json declares downloadable artifacts (URL, SHA256, archive name, provided directories). scripts/fetch_artifacts.py downloads one, verifies the checksum, unpacks into a staging directory and moves it into place, refusing to replace a non-empty destination without --force and rejecting archives with traversal, links or members outside the declared directories. Standard library only; Google Drive share links are rewritten to direct downloads. scripts/package_artifacts.py builds the reproducible tarball and prints the checksum for the manifest. The llm-aacl2026 entry keeps explicit placeholders for URL and SHA256 until the upload exists, and the downloader reports that instead of guessing. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA (cherry picked from commit 5c74683ebfaeba81268a1bcb665a3a8233ab9d82) * chore(artifacts): publish the llm-aacl2026 download URL and checksum The paper's LLM artifact is uploaded, so the manifest placeholders become the file's Google Drive share link and the SHA256 of the packaged archive. A real download through fetch_artifacts.py verified the checksum and restored 6140 files identical to the source tree. The manifest test now accepts either a fully published entry or an explicit placeholder, never a partial one. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA (cherry picked from commit 042f6d8d0460cc31cc0724cccf0d75a84b797fda) * refactor(results): move the B200 summary CSVs to results/B200/csv Path-only migration ahead of hardware-scoped results. All 90 committed CSVs (45 operators x {default, autotune}) are B200 measurements, so they move from results/csv/ to results/B200/csv/. Pure rename: no file is regenerated, reformatted or newline-normalised (56 of them use CRLF). Committed with --no-verify so the whitespace hooks cannot touch the data. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA * feat(results): scope results by hardware, results/<gpu>/... TileBench will hold results for several GPUs (B200 today; GH200 and an AMD part to follow), and a global results/{csv,logs,figures,aggregate,runs} lets the same operator overwrite itself across machines. Everything is now results/<gpu>/{csv,logs,figures,aggregate,runs}; only csv/ is tracked. - tilebench/paths.py: results_root / results_{csv,logs,figures,aggregate, runs}_dir and hardware_label(), which accepts any single safe path component. There is no device list: a new GPU needs a label, not a patch. - run_bench.py, run_bench_all.py, visualize.py, plot_sweep_max.py, aggregate_results.py, ncu_catalogue.py take --gpu. It is required and has no default, so a run on another machine cannot land in B200's namespace. run_bench.py prints the namespace next to the detected device and warns when the label does not appear in the device name. - A GPU CSV holds measurements from that GPU only. The TileLang merge is unchanged but confined to results/<gpu>/csv/. NKI (AWS Trainium) can no longer be merged into, or run together with, GPU backends: it is selected alone, without --gpu, and writes to outputs/nki/. Its profiler artifacts move there from results/logs/ too. The default backend set, and 'all', is triton, cutile, tilelang. - plot_sweep_max.py writes to results/<gpu>/figures/ by default so plotting another GPU cannot overwrite the README figure; pass --output for that. - .gitignore tracks results/<hardware>/csv/**/*.csv for any label and keeps logs, figures, aggregates (CSV too) and runs ignored. - archive_artifacts.sh --logs takes --gpu and snapshots results/<gpu>/logs/ only. Logs of every other GPU on the archive branch, and the legacy results/logs/ and benchmarks/llm_generated/ snapshots, are carried forward. run_bench.py archives the logs of its own --gpu. - Tests: path validation, per-GPU defaults for every script, no collision between GPUs, TileLang merge isolation, NKI isolation, .gitignore rules, archive behaviour, and the 90 B200 CSVs checked against SHA256 recorded before the move. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA * fix(results): keep NKI results inside the campaign namespace The previous commit isolated NKI under outputs/nki/. That is reverted: NKI keeps its original result organisation, moved into the hardware-scoped layout. `run_bench.py --gpu B200 --tile-language nki` is allowed again; --gpu names the result namespace of the campaign and does not claim that NKI ran on that GPU. NKI still runs through its own path on AWS Trainium, and engine and kernel execution are unchanged. - The NKI merge into results/<gpu>/csv/<op>_{default,autotune}.csv is restored as it is on main: torch_nki_ms, nki_ms and speedup_nki are appended as measured, with no B200 drift scaling, and speedup_nki = torch_nki_ms / nki_ms, never torch_ms / nki_ms. - NKI timing and autotune JSON go to results/<gpu>/logs/ like every other log. The Neuron profiling artifacts move from results/logs/ to results/<gpu>/logs/nki_profiles/ and nki_neff_manifest.jsonl: the engine takes logs_dir and hands both paths to the orchestrator, which no longer has a default location. - --gpu is required again for every run. NKI still only runs when named: the default selection and 'all' stay triton, cutile, tilelang. run_bench_all.py goes back to the engine's own backend selection. - archive_artifacts.sh --logs --gpu <gpu> snapshots the whole logs/ tree, NKI profiles included. There is no separate NKI archive path. - Tests: NKI merges into the B200 CSV, its columns survive a later TileLang merge, frozen GPU columns and the GH200 CSV stay untouched, speedup_nki uses torch_nki_ms, NKI logs and profiles land under results/<gpu>/logs/ and are archived with it. The B200 guard now accepts the NKI extension: the manifest records the pre-move SHA256 and a digest of the frozen columns, so a file is either byte-identical or extended with its GPU columns intact. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA * refactor(profiling): move the B200 NCU metadata to metadata/B200/ Path-only move ahead of hardware-scoped profiling metadata. The catalogue (sweep-max cases and autotune winners) and the probed kernel counts were recorded on B200, so they move from tilebench/profiling/ to tilebench/profiling/metadata/B200/. No content change. Committed with --no-verify: neither file ends in a newline, and the end-of-file hook would otherwise rewrite them. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA * feat: multi-hardware consistency for logs, NCU metadata and the archive Four places still assumed a single GPU or a single run per operator. Raw JSON. results/<gpu>/logs/ held one <op>_results.json and one <op>_autotune.json, so a default run, an autotune run, a TileLang-only run and an NKI run overwrote each other. The default names now carry the mode and the backend selection: <op>_<mode>_<backends>.json, built only by tilebench.paths.timing_log_path / autotune_log_path. tilebench/backends.py owns the backend names, their canonical order and the --tile-language parser, so `cutile,triton` and `triton,cutile` are one tag. Explicit --output / --autotune-log still win. With one file per selection the autotune-log merge is gone: it existed to stop a TileLang run clobbering the Triton/cuTile winners, and would now refuse a rerun with other cases. visualize.py names the run it reads with --mode and --tile-language, and ncu_catalogue.py reads exactly <op>_autotune_<backends>.json (default triton,cutile; both are required). Nothing is chosen by glob or mtime. NCU metadata. The catalogue and the kernel counts were one global pair, although sweep-max cases, autotune winners, launch counts and kernel names differ between GPUs. They are now tilebench/profiling/metadata/<gpu>/{ncu_catalogue,kernel_counts}.json (the B200 files moved in the previous commit), reached through ncu_catalogue_path / kernel_counts_path. NCU_CATALOGUE and KERNEL_COUNTS are removed. ncu_catalogue, probe_kernel_count, ncu_one, ncu_driver, ncu_writeup, hf_upload and plot_sweep_max take --gpu, required and without a default. A GPU without metadata is an error that names the command to produce it; another GPU's files are never used. Reports go to outputs/ncu/<gpu>/, and hf_upload uploads under ncu_report_main/<gpu>/ so it cannot overwrite the released B200 reports. package-data covers the nested directory, checked by a non-editable install. Archive. main still tracks the legacy results/logs/, so reading main's tree would bring it back on every run. It is now dropped from the base tree, and results/logs/ on an archive tip is re-rooted at results/B200/logs/ byte for byte. A migrated file that meets a different file at its destination, already archived or in the working directory, aborts the run with nothing committed. The logs of other GPUs and both LLM paths are carried forward as before. The summary CSVs come from main's tree and are never copied. Portability. The profiling shell scripts located the repository through a hard-coded /projects/... path, a user-specific conda path and QOS, and PYTHONPATH=. They resolve it from their own location (batch2.sbatch from the submit directory) and work from any directory. run_batch.sh no longer renames or copies results: run_bench.py owns the names, and the <op>_summary.csv it copied has not existed for a while. run_bench.py starts the archive script inside the repository whatever the caller's directory. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA * docs(readme): camera-ready README - New tagline (accelerator framework; cuTile on CUDA 13.2, Triton, TileLang, AWS Neuron NKI, PyTorch); NKI added to the backend list. - Figures under assets/: the paper's overview diagram and a new sweep-max latency chart (scripts/plot_sweep_max.py; PyTorch / Triton / cuTile at the NCU-catalogue case of each operator). - Software Versions section with the stack the results were measured on. - OPERATOR_AUTHORING_GUIDE.md merged in as a section; the duplicated Operator Config / Autotune / Adding a New Operator sections are dropped. - Stale content brought up to date: project structure, 2x-L2 eviction size, CLI flags (--autotune, --tile-language, --gpu), Triton and cuTile examples (best_config / CutileAutotuner, no global state), peaks read from data/peak_performance/<GPU>.json, FP8 status re-verified on this stack. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA (cherry picked from commit 24e28f5) * docs(readme): corrected overview figure, version note, auto-archive, skills/ - assets/overview.png: corrected figure ("Verify"). - Software Versions: the TileBench paper was measured with cuda-tile 1.3.0; the TileBench++ work (TileLang, NKI) uses 1.5.0. - Document the automatic raw-log archive, --no-archive, and skills/. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA (cherry picked from commit 483bc3f) * docs(readme): restructure around the paper results Lead with the paper's headline results (pairwise win/loss matrix, suite geomeans, autotune gains) and the sweep-max latency figure, then the benchmark overview, backend scope, paper vs. development environments, quick start, evaluation methodology, the LLM track, and a collapsible developer reference (CLI tables and the operator authoring guide). Checked against the paper sources and the recorded CSVs: the win/loss matrix (37/8, 36/9, 33/12; 14 operators within 5%) recomputes exactly from results/csv, every operator sweeps 20 cases per dtype, and the section, appendix and table references resolve. The quick-start install no longer lists cuda-tile-experimental: it is not on PyPI and nothing in the repository imports it. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA (cherry picked from commit 2a85c11) * docs: add developer guide (cherry picked from commit 14198b3) * docs: rewrite README for benchmark repository (cherry picked from commit 29ca688) * docs: add TileBench icon (cherry picked from commit 2a807b9) * docs: add TileBench icon to README (cherry picked from commit 4c9cb6c) * Clean up README.md formatting Removed extra blank lines in README.md (cherry picked from commit ddad492) * docs: describe LLM artifacts as local output with a one-command restore Separate the three things that used to read as one: the version-controlled pipeline under tilebench/llm_codegen/, the git-ignored local output under tilebench/benchmarks/llm_generated/, and the paper's frozen snapshot restored with fetch_artifacts.py. The download link is marked as not yet published. Recorded Results now lists where each kind of generated artifact is kept, and the developer guide documents archive_artifacts.sh and the publishing flow. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA * docs(readme): link the published LLM artifact Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA * docs(readme): resource links with icons, and the published NCU reports Add a resource row under the description, in the spirit of KernelBench's arXiv | blog | dataset line: NCU reports on Hugging Face, the LLM artifact on Google Drive, the raw-log archive branch and the developer guide, each led by its service's icon. The same icons mark those links where they appear in the body and in the Recorded Results table, which gains a row for the Nsight Compute reports. The 220 reports (45 operators, Triton and cuTile, every profiled dtype, 11.8 GB) are already public at huggingface.co/datasets/bcui2/NCU_report; the README now says so, and that downloading needs a free Hugging Face login because the dataset is gated with automatic approval. Icons are vendored under assets/icons/ rather than hot-linked: the Hugging Face logo from its brand-assets repository (viewBox cropped to the glyph), the Google Drive product icon from Wikimedia Commons, the GitHub mark from Simple Icons and the book glyph from Octicons, the last two in a neutral grey that stays visible in GitHub's dark theme. Each icon sits inside its link's text, so GitHub does not wrap it in a link to the image file. Also fixes the tagline, which followed </div> without a blank line and was therefore rendered with its literal ** markers. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA * docs(readme): TileBench logo, transparent overview, quieter inline links - Put the TileBench logo at the top as the page title. The supplied artwork is cropped to its content (top and bottom margins removed). The light variant has the page background removed; the dark variant keeps the original cream card so the black wordmark stays readable. - Replace the overview figure with the transparent-background version. It is pixel-identical to the previous one when flattened on white. - Keep icons in the top resource bar and the Recorded Results table only; inline body links are plain Markdown. Section anchors are unchanged. - NCU wording: "The released artifact contains 220 raw Nsight Compute reports." - State that tilebench/benchmarks/llm_generated/ is the Git-ignored runtime output path, that a fresh clone does not ship the 6140 campaign files, and that fetch_artifacts.py restores them automatically. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA * Revise README overview for clarity and detail Updated the overview section to provide a clearer description of TileBench and its supported backends. * docs(readme): add an icon to every section heading Emoji prefix on the 14 `##` headings, in the style of the SWE-bench README. GitHub drops the emoji from the heading slug but keeps the following space, so the 12 table-of-contents links move from `#name` to `#-name`. No other content changes. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA * docs: hardware-scoped results, results/<gpu>/..., and the --gpu label Follows the results layout change on the infra branch. Every command in the README and the developer guide passes --gpu B200, and the result paths read results/<hardware>/{csv,logs,figures,aggregate,runs}/. The text states that this is a general layout, that results/B200/csv/ holds the paper results committed today, and that other hardware sits beside it without a code change. The developer guide documents the label rules, the one-namespace- one-hardware rule, TileLang merge scope, NKI output under outputs/nki/, and the per-GPU archive command. Headline B200 numbers are unchanged. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA * docs: NKI results are kept with the B200 campaign, as cross-hardware data Follows the correction on the infra branch. The previous text said NKI was kept out of the GPU namespaces, under outputs/nki/. NKI keeps its original organisation instead: its columns (torch_nki_ms, nki_ms, speedup_nki) are merged into results/B200/csv/, and its logs and profiles live under results/B200/logs/. The README and the developer guide now say that these are measurements taken on AWS Trainium, stored beside the B200 results for a unified per-operator record, that speedup_nki = torch_nki_ms / nki_ms and never torch_ms / nki_ms, and that --gpu names the campaign rather than the device NKI ran on. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA * docs: per-hardware NCU metadata, collision-free raw JSON, migrated archive logs Follows the multi-hardware consistency work on the infra branch. - Profiling metadata is documented as tilebench/profiling/metadata/<hardware>/{ncu_catalogue,kernel_counts}.json, with B200 as the directory committed today and other hardware beside it. NCU reports are written to outputs/ncu/<hardware>/. The 220 released reports are described as the paper's B200 profiles; the count is unchanged. - Raw JSON names carry the mode and the backend selection. The README shows the two files its quick start produces and passes --tile-language to visualize.py, which now names the run it reads. The developer guide adds the naming rule, --mode, the ncu_catalogue log contract and the --gpu commands of the NCU tools. - The paper's B200 raw logs are stated to live under results/B200/logs/ on the archive branch. The sentence describing results/logs/ as their current location is gone. The Backend Support table is unchanged, TileLang "N/A" included. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA * Clean profiling package layout tilebench/ is the installable Python source. tilebench/profiling/ still held cluster launch scripts and measured data, which are neither. - Remove batch1_launch.sh, batch2.sbatch, rerun_timeout.sh, run_batch.sh and supplemental_launch.sh. They are operational scripts of one campaign on one cluster; no source file, script or test of the framework depends on them. They are kept on the archive branch, not moved to scripts/. - Move the NCU metadata out of the package. ncu_catalogue.json (autotune winners) and kernel_counts.json (probed launch counts and kernel names) are measured on one GPU and written by ncu_catalogue.py / probe_kernel_count.py: generated experiment data, not a package resource. No benchmark code path reads them; only the NCU tools do. They now live under the Git-ignored outputs/profiling/<gpu>/, next to the NCU reports, still one directory per GPU and still with no fallback to another GPU. The committed B200 files leave the tree and are kept on the archive branch. - plot_sweep_max.py no longer needs that metadata: the sweep-max case comes from the operator's config.yaml through the same rule the catalogue uses (ncu_catalogue.sweep_max_cases), so the README figure is reproducible from the configs and the committed CSVs alone. - pyproject.toml: drop the tilebench.profiling package-data entry. Wheel and sdist now carry only .py files for that package. - Tests: drop the two tests that only exercised the removed shell scripts, and the SHA256 fingerprints of the B200 CSVs and of the B200 metadata (tests/data/ is gone). They proved one migration, not a software contract. The CSV schema test, the namespace and isolation tests and the NKI merge tests stay. A packaging test now asserts that a regular install ships no shell script and no data under tilebench/profiling/. - archive_artifacts.sh: carry forward outputs/profiling/ and any launch script kept on the archive tip, and keep the tip's own tree as the base when it already contains origin/main, so code merged into the archive from a branch ahead of main is not rolled back by the next archive run. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA * docs: NCU metadata is generated data under outputs/profiling/<hardware>/ Follows "Clean profiling package layout" on repo-cleanup. The README and the developer guide described the NCU catalogue and kernel counts as files committed under tilebench/profiling/metadata/. They are generated, hardware-specific data and now live under the Git-ignored outputs/profiling/<hardware>/; nothing is committed for any GPU, and tilebench/profiling/ is described as source only. The Recorded Results table gains one row for this metadata. The developer guide notes that plot_sweep_max.py needs only the configs and the committed CSVs, and what the archive script carries forward. No removed launch script was referenced. The Backend Support table is unchanged. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA * Separate profiling tools from the profiling library tilebench/profiling/ still mixed three kinds of Python file. "All .py" is not the same as "package source", so each of the 17 files was classified by what it does and who imports it. Library, stays in the package (3): __init__.py ncu_kernel_select.py kernel selection, capture validation and metadata loading, imported by four tools ncu_catalogue.py sweep-max cases and catalogue entries; imported by plot_sweep_max.py. Its argparse main and sys.path hack are gone; write_catalogue() is the entry point. Command-line tools, moved to scripts/ (8): scripts/aggregate_results.py aggregates result CSVs, so it sits with visualize.py, not with NCU scripts/profiling/ncu_catalogue.py thin CLI over the library scripts/profiling/{ncu_driver,ncu_one,ncu_writeup,probe_kernel_count, hf_upload}.py had no importer outside the tests scripts/profiling/ncu_generic_harness.py never imported: it is the process NCU profiles, started by path They bootstrap sys.path like the other scripts and run from any directory without PYTHONPATH. The drivers locate the harness next to themselves instead of inside the package. The report path rule both drivers duplicated is now tilebench.paths.ncu_report_path. Campaign-specific, removed from the public tree (7), kept on the archive: probe_hung_cfgs.py and six ncu_harness_*.py. Hard-coded operators, cases and configs, code that runs at import, the pre-profiler-range methodology the generic harness replaced, and no reference from any driver, test or document. NCU behaviour, metadata and output paths are unchanged: a real capture of mul2 on a B200, started from /tmp with the relocated tools, validated 1/1 launches. archive_artifacts.sh now carries forward whatever the archive tip keeps under tilebench/profiling/ that the base tree dropped, not only *.sh/*.sbatch. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA * docs: profiling tools are scripts, tilebench/profiling is the library Follows "Separate profiling tools from the profiling library" on repo-cleanup. Every profiling command changes from `python -m tilebench.profiling.<tool>` to `python scripts/profiling/<tool>.py`, and aggregate_results becomes `python scripts/aggregate_results.py`. The developer guide replaces "tilebench/profiling/ holds Python source only" with the actual rule: a module is in the package only if other code imports it, programs live under scripts/, and campaign-specific files and measured data belong in neither. The README project structure lists scripts/profiling/. The Backend Support table is unchanged. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA * Decouple run_bench.py from the archive workflow A benchmark run only produces files. By default run_bench.py also fetched origin/main and the 7,000-file archive branch, and committed a snapshot of the raw logs to a local branch of the user's clone; the README had to tell readers to pass --no-archive on every command. Backing results up is a maintainer task that targets one dated archive branch, not a runtime feature. - run_bench.py: remove _archive_logs(), the --no-archive flag and the call after the run, with the subprocess and REPO_ROOT imports only they used. There is no opt-in replacement and no compatibility no-op. - Remove scripts/archive_artifacts.sh, scripts/archive_logs.sh and tests/test_archive_scripts.py. They are kept on the archive branch, where archive_artifacts.sh remains the maintenance tool of that branch. - results/<hardware>/logs/ stays Git-ignored: a run writes its logs there and performs no Git operation. - Tests: drop the archive hook from the run_bench fixture and the test of the hook itself; remove --no-archive from every call. A new test asserts that a run starts no process and that --no-archive and --archive are rejected. Namespace isolation, NKI merge, path layout and raw JSON collision tests are unchanged. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA * docs: running a benchmark archives nothing Follows "Decouple run_bench.py from the archive workflow" on repo-cleanup. --no-archive is gone from the three quick-start commands and from the run_bench.py argument table, the statement that raw logs are snapshotted after each run is gone, and so are the maintainer archive command and the "Archiving artifacts" section of the developer guide. What remains is factual: the paper's frozen raw logs, LLM trajectories and B200 NCU metadata are preserved on archive/raw-logs-2026-09-18, and a benchmark run performs no Git operation. The Backend Support table is unchanged. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA * tests: the no-side-effect test names no removed flag Keep the assertion that matters, that run_bench.py starts no process, and drop the check that two archive flags are rejected: argparse rejects any unknown flag, and the public tree should not name an option it never had. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA * refactor: move problems/ to tilebench/problems and rename llm_codegen to llm Path-only. tilebench/llm_codegen/problems/ becomes tilebench/problems/, a sibling of benchmarks/, core/ and data/, and tilebench/llm_codegen/ becomes tilebench/llm/. No file content changes; the next commit updates the references. Committed with --no-verify: the task descriptions are prompt inputs and must stay byte-identical, so the whitespace hooks must not touch them. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA * Point the code at tilebench/problems and tilebench/llm Follows the path-only move. tilebench.paths gains PROBLEMS_ROOT and renames LLM_CODEGEN_ROOT to LLM_ROOT; prompt_builder.py and generate_descriptions.py read the task descriptions through PROBLEMS_ROOT instead of a directory next to themselves. Module references become tilebench.llm, run as `python -m tilebench.llm.generate`. tilebench/problems/ holds only Markdown, so pyproject.toml declares it as data of the top-level package, and migrate_pr_paths.py maps main's benchmarks/problems/current/ and tools/llm_codegen/ to the new locations. The 45 rebuilt initial prompts are byte-identical to the ones built before the move. A wheel and an sdist carry the 45 task descriptions and the llm package; tests assert that every operator has a description under PROBLEMS_ROOT and that a regular install ships them. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA * docs: the LLM pipeline is tilebench/llm, the task descriptions tilebench/problems Follows the move on repo-cleanup: tilebench/llm_codegen/ is now tilebench/llm/, and its problems/ directory is tilebench/problems/, beside benchmarks/, core/ and data/. The README section, the project structure and the developer guide use the new paths; the guide names PROBLEMS_ROOT and the `python -m tilebench.llm.generate` entry point. The Backend Support table is unchanged. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01EdNBRsQvET5PUJ3Nw1hVCA --------- Co-authored-by: Claude Fable 5.1 <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
2 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.
NKI (AWS Trainium) implementation of mul2, split out of the consolidated NKI branch
cecilia/feature/nki-vector-add(nki-all-operators, #259) so each operator can be reviewed independently. Supersedes #185 (older per-operator branch: legacyneuronxcc.nkiimports; this is the migratedimport nkiversion).Files: A benchmarks/operators/mul2/impl_nki.py
Status: imports and exposes run()/get_last_config() on trn2 (nki 0.6.0); not individually re-benchmarked in this split
Implementation by @Cecilia123li. Timing/identity infrastructure: #261; Trainium peak/roofline infra: #262.
🤖 Generated with Claude Code
https://claude.ai/code/session_012Q38kGmXvyoeM1qtCbheSL
Autotune (b383f26)
block_size(free-dim tile per 128-partition block)BLOCK_SIZEautotune=Falsekeeps the previous constants (default numbers unchanged). Validation on trn2, case 0 (default run + autotune code path with the candidate timer stubbed — no sweep;--autotuneruns a real sweep):