From fe447c43d1c08c48138d01d3bc0cd2d1501108d2 Mon Sep 17 00:00:00 2001 From: Brandon Ros Date: Mon, 14 Sep 2026 17:52:49 -0400 Subject: [PATCH 1/2] Add workload sweeps, per-module comparisons, and GPU experiment runners --- .github/workflows/optimization_wave2.yml | 68 ++++++++ .github/workflows/ptx_export.yml | 102 ++++++++++-- Cargo.lock | 7 + Cargo.toml | 1 + examples/ptx_export/check_module_cleanup.py | 100 ++++++++++++ examples/ptx_export/check_workloads.py | 76 +++++++++ examples/ptx_export/diagnose_wave2_nvvm.py | 61 +++++++ .../representative-kernels/Cargo.toml | 11 ++ .../representative-kernels/src/lib.rs | 60 +++++++ examples/ptx_export/run_guarded_select.py | 92 +++++++++++ examples/ptx_export/run_wave2.py | 102 ++++++++++++ examples/ptx_export/run_wave2_gpu.py | 152 ++++++++++++++++++ examples/ptx_export/test_module_cleanup.py | 28 ++++ 13 files changed, 851 insertions(+), 9 deletions(-) create mode 100644 .github/workflows/optimization_wave2.yml create mode 100644 examples/ptx_export/check_module_cleanup.py create mode 100644 examples/ptx_export/check_workloads.py create mode 100644 examples/ptx_export/diagnose_wave2_nvvm.py create mode 100644 examples/ptx_export/representative-kernels/Cargo.toml create mode 100644 examples/ptx_export/representative-kernels/src/lib.rs create mode 100644 examples/ptx_export/run_guarded_select.py create mode 100644 examples/ptx_export/run_wave2.py create mode 100644 examples/ptx_export/run_wave2_gpu.py create mode 100644 examples/ptx_export/test_module_cleanup.py diff --git a/.github/workflows/optimization_wave2.yml b/.github/workflows/optimization_wave2.yml new file mode 100644 index 00000000..98e31250 --- /dev/null +++ b/.github/workflows/optimization_wave2.yml @@ -0,0 +1,68 @@ +name: LLVM 21 optimization wave 2 + +on: + workflow_dispatch: + inputs: + workload: + description: Workload to repeat (all runs the complete matrix) + type: choice + default: all + options: [all, small, representative, solana, bitcoin, ethereum, shallenge, self_test] + push: + branches: [poc/portable-ptx-export] + paths: + - '.github/workflows/optimization_wave2.yml' + - 'examples/ptx_export/run_wave2.py' + - 'examples/ptx_export/representative-kernels/**' + +permissions: + contents: read + +jobs: + compare: + name: Compare ${{ matrix.workload }} + runs-on: ubuntu-24.04 + timeout-minutes: ${{ matrix.workload == 'self_test' && 90 || 45 }} + strategy: + fail-fast: false + matrix: + workload: ${{ fromJSON(inputs.workload != '' && inputs.workload != 'all' && format('["{0}"]', inputs.workload) || '["small","representative","solana","bitcoin","ethereum","shallenge","self_test"]') }} + steps: + - uses: actions/checkout@v4 + - uses: DeterminateSystems/nix-installer-action@main + - uses: DeterminateSystems/magic-nix-cache-action@v14 + with: + use-flakehub: false + - uses: actions/cache/restore@v4 + with: + path: | + ~/.cargo/registry + ~/.cargo/git + target/ + key: wave2-cargo-${{ runner.os }}-llvm21-${{ github.sha }} + restore-keys: | + ptx-cargo-v1-${{ runner.os }}-${{ runner.arch }}-llvm21- + - name: Build the checked-out backend + run: | + nix develop .#v21 --command cargo build -p rustc_codegen_nvvm --features llvm21 --target-dir target/cuda-builder-codegen + rm -rf target/nvptx64-nvidia-cuda + - name: Checkout pinned workload + uses: actions/checkout@v4 + with: + repository: brandonros/vanity-miner-rs + ref: 9791234249fc8cb762c296c4fda4503d2686ff77 + path: workloads/vanity-miner + - name: Compile and compare + run: | + mkdir -p artifacts + git rev-parse HEAD > artifacts/source-commit.txt + sha256sum target/cuda-builder-codegen/debug/librustc_codegen_nvvm.so > artifacts/backend-sha256.txt + nix develop .#v21 --command rustc -Vv > artifacts/rustc-version.txt + cp flake.lock Cargo.lock rust-toolchain.toml artifacts/ + nix develop .#v21 --command python3 examples/ptx_export/run_wave2.py ${{ matrix.workload }} --out artifacts/${{ matrix.workload }} --miner workloads/vanity-miner + - uses: actions/upload-artifact@v4 + if: ${{ !cancelled() }} + with: + name: wave2-${{ matrix.workload }} + path: artifacts/ + if-no-files-found: error diff --git a/.github/workflows/ptx_export.yml b/.github/workflows/ptx_export.yml index 4b57305b..527d7546 100644 --- a/.github/workflows/ptx_export.yml +++ b/.github/workflows/ptx_export.yml @@ -4,6 +4,19 @@ on: pull_request: branches: [main, experiment/cuda13.3-llvm21, "stack/**"] workflow_dispatch: + inputs: + cleanup_experiment: + description: Replay LLVM 21 handoff IR with opt-in cleanup passes + type: boolean + default: false + optimization_sweep: + description: Run the extended CFG, DCE, memory and inlining experiments + type: boolean + default: false + mining_workload: + description: Compile pinned Solana mining source and replay memory cleanup + type: boolean + default: false push: branches: [poc/portable-ptx-export] @@ -17,23 +30,94 @@ jobs: steps: - uses: actions/checkout@v4 - uses: DeterminateSystems/nix-installer-action@main + - uses: DeterminateSystems/magic-nix-cache-action@v14 + with: + use-flakehub: false + - name: Restore Cargo cache + id: cargo-cache + uses: actions/cache/restore@v4 + with: + # Include all nested builds, notably target/cuda-builder-codegen. + path: | + ~/.cargo/registry + ~/.cargo/git + target/ + key: ptx-cargo-v1-${{ runner.os }}-${{ runner.arch }}-llvm21-${{ hashFiles('flake.nix', 'flake.lock', 'rust-toolchain.toml', '**/Cargo.toml', '**/Cargo.lock', '.cargo/**', '.github/workflows/ptx_export.yml') }}-${{ github.sha }} + restore-keys: | + ptx-cargo-v1-${{ runner.os }}-${{ runner.arch }}-llvm21-${{ hashFiles('flake.nix', 'flake.lock', 'rust-toolchain.toml', '**/Cargo.toml', '**/Cargo.lock', '.cargo/**', '.github/workflows/ptx_export.yml') }}- + - name: Test codegen inventory parser + run: python3 -m unittest discover -s examples/ptx_export -p 'test_*.py' + - name: Rebuild the checked-out backend and regenerate device evidence + run: | + # cuda_builder searches for an existing backend before trying Cargo. + # A restored .so alone therefore does not establish source freshness. + nix develop .#v21 --command cargo build -p rustc_codegen_nvvm --features llvm21 --target-dir target/cuda-builder-codegen + # Device Cargo fingerprints do not track changes to the backend .so; + # cached builds also omit the IR side outputs required by this job. + # Retain the expensive host/backend dependency cache. + rm -rf target/nvptx64-nvidia-cuda + - name: Check guarded-select source semantics on CPU + run: | + nix develop .#v21 --command rustc --edition=2024 --test examples/ptx_export/kernels/src/guarded_select.rs -o /tmp/guarded-select-debug + /tmp/guarded-select-debug + nix develop .#v21 --command rustc --edition=2024 -O --test examples/ptx_export/kernels/src/guarded_select.rs -o /tmp/guarded-select-release + /tmp/guarded-select-release - name: Compile Rust kernels without a GPU - run: nix develop .#v21 --command cargo run -p ptx_export --features llvm21 -- artifacts/ptx + id: export + run: | + mkdir -p artifacts/ptx + printf '%s\n' 'nix develop .#v21 --command cargo run -vv -p ptx_export --features llvm21 -- artifacts/ptx none' > artifacts/ptx/build-command.txt + nix develop .#v21 --command cargo run -vv -p ptx_export --features llvm21 -- artifacts/ptx none 2>&1 | tee artifacts/ptx/build.log + - name: Verify LLVM 21 handoff IR + run: nix develop .#v21 --command opt-21 -passes=verify -disable-output artifacts/ptx/final-module.ll - name: Record provenance run: | git rev-parse HEAD > artifacts/ptx/source-commit.txt nix develop .#v21 --command rustc -Vv > artifacts/ptx/rustc-version.txt - sha256sum artifacts/ptx/rust_kernels.ptx > artifacts/ptx/SHA256SUMS - - name: Check Python regressions - run: python3 -B -m unittest discover -s examples/ptx_export -p 'test_*.py' - - name: Check guarded-select semantics - run: | - nix develop .#v21 --command rustc --edition=2024 --test examples/ptx_export/kernels/src/guarded_select.rs -o /tmp/guarded-select - /tmp/guarded-select - - name: Check default DCE retention + nix develop .#v21 --command nvcc --version > artifacts/ptx/nvcc-version.txt + cp flake.lock rust-toolchain.toml Cargo.lock artifacts/ptx/ + sha256sum target/cuda-builder-codegen/debug/librustc_codegen_nvvm.so > artifacts/ptx/backend-sha256.txt + # Keep per-crate LLVM IR as well as the linked final-module.ll. + mkdir -p artifacts/ptx/llvm-ir + find target/nvptx64-nvidia-cuda -name '*.ll' -exec cp --parents '{}' artifacts/ptx/llvm-ir/ \; + - name: Assemble and inspect NVIDIA machine code without a GPU + run: nix develop .#v21 --command python3 examples/ptx_export/inspect_codegen.py artifacts/ptx + - name: Verify default DCE retention and explicit disable run: nix develop .#v21 --command python3 examples/ptx_export/check_default_dce.py artifacts/ptx + - name: Investigate pre-NVVM cleanup (opt-in) + if: ${{ github.event_name == 'workflow_dispatch' && (inputs.cleanup_experiment || inputs.optimization_sweep) }} + run: nix develop .#v21 --command python3 examples/ptx_export/replay_cleanup.py artifacts/ptx ${{ inputs.optimization_sweep && '--extended' || '' }} + - name: Validate integrated cleanup against replay (opt-in) + if: ${{ github.event_name == 'workflow_dispatch' && (inputs.cleanup_experiment || inputs.optimization_sweep) }} + run: nix develop .#v21 --command python3 examples/ptx_export/check_integrated_cleanup.py artifacts/ptx + - name: Validate per-module cleanup (extended sweep) + if: ${{ !cancelled() && steps.export.outcome == 'success' && github.event_name == 'workflow_dispatch' && inputs.optimization_sweep }} + run: nix develop .#v21 --command python3 examples/ptx_export/check_module_cleanup.py artifacts/ptx + - name: Validate size-oriented builds (extended sweep) + if: ${{ !cancelled() && steps.export.outcome == 'success' && github.event_name == 'workflow_dispatch' && inputs.optimization_sweep }} + run: nix develop .#v21 --command python3 examples/ptx_export/check_workloads.py artifacts/ptx + - name: Checkout pinned mining workload + if: ${{ !cancelled() && steps.export.outcome == 'success' && github.event_name == 'workflow_dispatch' && inputs.mining_workload }} + uses: actions/checkout@v4 + with: + repository: brandonros/vanity-miner-rs + ref: 9791234249fc8cb762c296c4fda4503d2686ff77 + path: workloads/vanity-miner + - name: Compare cleanup on Solana mining kernel + if: ${{ !cancelled() && steps.export.outcome == 'success' && github.event_name == 'workflow_dispatch' && inputs.mining_workload }} + run: nix develop .#v21 --command python3 examples/ptx_export/check_workloads.py artifacts/ptx --miner workloads/vanity-miner - uses: actions/upload-artifact@v4 + if: ${{ !cancelled() }} with: name: rust-ptx path: artifacts/ptx/ if-no-files-found: error + - name: Save Cargo cache, including after compilation failures + if: ${{ !cancelled() && steps.cargo-cache.outcome == 'success' && steps.cargo-cache.outputs.cache-hit != 'true' }} + uses: actions/cache/save@v4 + with: + path: | + ~/.cargo/registry + ~/.cargo/git + target/ + key: ${{ steps.cargo-cache.outputs.cache-primary-key }} diff --git a/Cargo.lock b/Cargo.lock index a4fa73c7..97727a62 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -2592,6 +2592,13 @@ dependencies = [ "sha2", ] +[[package]] +name = "ptx-representative-kernels" +version = "0.1.0" +dependencies = [ + "cuda_std", +] + [[package]] name = "ptx-retention-kernels" version = "0.1.0" diff --git a/Cargo.toml b/Cargo.toml index d872f361..fac7da6b 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -54,6 +54,7 @@ members = [ "examples/ptx_export", "examples/ptx_export/kernels", "examples/ptx_export/retention-kernels", + "examples/ptx_export/representative-kernels", "samples/introduction/async_api", "samples/introduction/async_api/kernels", diff --git a/examples/ptx_export/check_module_cleanup.py b/examples/ptx_export/check_module_cleanup.py new file mode 100644 index 00000000..42abe4b5 --- /dev/null +++ b/examples/ptx_export/check_module_cleanup.py @@ -0,0 +1,100 @@ +#!/usr/bin/env python3 +"""Compare opt-in per-codegen-unit cleanup with independent LLVM replay.""" +import argparse +import hashlib +import json +import re +from pathlib import Path +import subprocess +import sys +from optimization_pipelines import SCALAR + + +def canonical_phi(line): + # PHI incoming pairs are unordered. Keep each value attached to its block; + # this is applied only after LLVM has verified and stripped local names. + if ' = phi ' not in line: + return line + groups = [] + depth = 0 + quoted = escaped = False + start = 0 + for index, char in enumerate(line): + if not quoted: + if char == '[': + if depth == 0: start = index + depth += 1 + elif char == ']': + depth -= 1 + if depth == 0: + group = line[start:index + 1] + if re.search(r',\s*%[\w.$-]+\s*\]$', group): + groups.append((start, index + 1, group)) + if char == '"' and not escaped: quoted = not quoted + escaped = char == "\\" and not escaped + if len(groups) < 2: + return line + if any(line[a[1]:b[0]].strip() != ',' for a, b in zip(groups, groups[1:])): + return line + return line[:groups[0][0]] + ', '.join(sorted(g[2] for g in groups)) + line[groups[-1][1]:] + + +def canonical(ir): + # Ignore LLVM printer comments (including predecessor order), but preserve + # semicolons inside quoted identifiers, inline assembly and string constants. + lines = [] + for line in ir.splitlines(): + quoted = escaped = False + end = len(line) + for index, char in enumerate(line): + if char == ';' and not quoted: + end = index + break + if char == '"' and not escaped: + quoted = not quoted + escaped = char == "\\" and not escaped + text = line[:end].rstrip() + if text.strip(): lines.append(canonical_phi(text)) + return '\n'.join(lines) + + +def main(): + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument('artifacts', type=Path) + args = parser.parse_args() + root = args.artifacts.resolve() + results = [] + for mode in ('module-scalar', 'module-inline'): + out = root/'module-cleanup'/mode; out.mkdir(parents=True, exist_ok=True) + command = ['cargo', 'run', '-vv', '-p', 'ptx_export', '--features', 'llvm21', '--', str(out), mode] + (out/'compiler-command.json').write_text(json.dumps(command, indent=2)+'\n') + with (out/'build.log').open('w') as log: + subprocess.run(command, stdout=log, stderr=subprocess.STDOUT, check=True) + before = sorted((out/'per-module').glob('*.before.ll')) + if not before: raise RuntimeError('no per-module before/after evidence') + checked = [] + for path in before: + after = path.with_name(path.name.replace('.before.ll', '.after.ll')) + replay = path.with_name(path.name.replace('.before.ll', '.replay.ll')) + normalized = path.with_name(path.name.replace('.before.ll', '.normalized.ll')) + subprocess.run(['opt-21', '-passes='+SCALAR+',strip-nondebug,verify', '-verify-each', '-S', str(path), '-o', str(replay)], check=True) + subprocess.run(['opt-21', '-passes=strip-nondebug,verify', '-S', str(after), '-o', str(normalized)], check=True) + if canonical(replay.read_text()) != canonical(normalized.read_text()): + raise RuntimeError(f'per-module replay mismatch: {path.name}') + checked.append({'module':path.name.removesuffix('.before.ll'), + 'before_sha256':hashlib.sha256(path.read_bytes()).hexdigest(), + 'after_sha256':hashlib.sha256(after.read_bytes()).hexdigest()}) + if not any(x['module'].startswith('core') for x in checked): + raise RuntimeError('missing dependency module coverage') + subprocess.run(['opt-21', '-passes=verify', '-disable-output', str(out/'final-module.ll')], check=True) + subprocess.run([sys.executable, str(Path(__file__).with_name('inspect_codegen.py')), str(out)], check=True) + if mode == 'module-inline': + subprocess.run([sys.executable, str(Path(__file__).with_name('check_cleanup_ir.py')), + str(out/'final-module.ll'), '--out', str(out/'host-ir-check')], check=True) + results.append({'mode':mode, 'checked_modules':checked}) + (root/'module-cleanup'/'comparison.json').write_text(json.dumps(results, indent=2)+'\n') + print(f'{mode}: {len(checked)} modules match standalone replay; final PTX assembled', flush=True) + + +if __name__ == '__main__': + main() diff --git a/examples/ptx_export/check_workloads.py b/examples/ptx_export/check_workloads.py new file mode 100644 index 00000000..ec77874e --- /dev/null +++ b/examples/ptx_export/check_workloads.py @@ -0,0 +1,76 @@ +#!/usr/bin/env python3 +"""Build size-oriented reproducers or pinned, unmodified Solana kernel sources.""" +import argparse +import json +from pathlib import Path +import re +import shutil +import subprocess +import sys +from replay_cleanup import normalized_functions + + +def main(): + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument('artifacts', type=Path) + parser.add_argument('--miner', type=Path) + args = parser.parse_args() + root = args.artifacts.resolve() + scripts = Path(__file__).resolve().parent + if args.miner: + miner = args.miner.resolve() + out = root/'mining-solana'; out.mkdir(parents=True, exist_ok=True) + commit = subprocess.check_output(['git', '-C', str(miner), 'rev-parse', 'HEAD'], text=True).strip() + if commit != '9791234249fc8cb762c296c4fda4503d2686ff77': + raise RuntimeError('unexpected mining workload revision') + manifest = miner/'kernels/Cargo.toml' + original = manifest.read_text() + shutil.copy2(manifest, out/'Cargo.toml.original') + shutil.copy2(miner/'kernels/Cargo.lock', out/'Cargo.lock.original') + # Use this backend's cuda_std while retaining the mining algorithm source. + replacement = 'cuda_std = { path = '+json.dumps(str(scripts.parents[1]/'crates/cuda_std'))+' }' + patched, count = re.subn(r'^cuda_std = \{ git = "https://github.com/brandonros/Rust-CUDA.git", rev = "2f4fd1d" \}$', replacement, original, flags=re.M) + if count != 1: raise RuntimeError('unexpected cuda_std dependency; review workload pin') + manifest.write_text(patched) + shutil.copy2(manifest, out/'Cargo.toml.patched') + (out/'workload.json').write_text(json.dumps({'repository':'brandonros/vanity-miner-rs', 'commit':commit, + 'features':['solana'], 'source_change':'cuda_std dependency path only; algorithm source unchanged'}, indent=2)+'\n') + builds = [('none', out, [str(miner/'kernels'), 'solana'])] + else: + builds = [(mode, root/'size-builds'/mode, []) for mode in ('size-s', 'size-z')] + for mode, out, extra in builds: + out.mkdir(parents=True, exist_ok=True) + command = ['cargo', 'run', '-vv', '-p', 'ptx_export', '--features', 'llvm21', '--', str(out), mode, *extra] + (out/'compiler-command.json').write_text(json.dumps(command, indent=2)+'\n') + with (out/'build.log').open('w') as log: + subprocess.run(command, stdout=log, stderr=subprocess.STDOUT, check=True) + subprocess.run(['opt-21', '-passes=verify', '-disable-output', str(out/'final-module.ll')], check=True) + subprocess.run([sys.executable, str(scripts/'inspect_codegen.py'), str(out)], check=True) + if not args.miner: + subprocess.run([sys.executable, str(scripts/'check_cleanup_ir.py'), str(out/'final-module.ll'), + '--out', str(out/'host-ir-check')], check=True) + if args.miner: + shutil.copy2(miner/'kernels/Cargo.lock', out/'Cargo.lock.resolved') + subprocess.run([sys.executable, str(scripts/'replay_cleanup.py'), str(out), '--extended', '--only', + 'baseline,dce-only,inline-only,inline-cleanup,memory-early-cse,memory-gvn,memory-stores,memory-combined'], check=True) + results = [] + for mode, replay in [('dce', 'dce-only'), ('inline-scalar', 'inline-only')]: + dest = out/'integrated-cleanup'/mode + dest.mkdir(parents=True, exist_ok=True) + command = ['cargo', 'run', '-vv', '-p', 'ptx_export', '--features', 'llvm21', '--', + str(dest), mode, str(miner/'kernels'), 'solana'] + (dest/'compiler-command.json').write_text(json.dumps(command, indent=2)+'\n') + with (dest/'build.log').open('w') as log: + subprocess.run(command, stdout=log, stderr=subprocess.STDOUT, check=True) + subprocess.run(['opt-21', '-passes=verify', '-disable-output', str(dest/'final-module.ll')], check=True) + actual = normalized_functions((dest/'rust_kernels.ptx').read_text()) + expected = normalized_functions((out/'cleanup-experiment'/replay/'rust_kernels.ptx').read_text()) + if 'kernel_find_solana_vanity_private_key' not in actual or actual != expected: + raise RuntimeError(f'Solana {mode}: integrated cleanup differs from replay') + subprocess.run([sys.executable, str(scripts/'inspect_codegen.py'), str(dest)], check=True) + results.append({'mode':mode, 'matches_replay_function_bodies':True}) + (out/'integrated-cleanup/comparison.json').write_text(json.dumps(results, indent=2)+'\n') + + +if __name__ == '__main__': + main() diff --git a/examples/ptx_export/diagnose_wave2_nvvm.py b/examples/ptx_export/diagnose_wave2_nvvm.py new file mode 100644 index 00000000..c640c53b --- /dev/null +++ b/examples/ptx_export/diagnose_wave2_nvvm.py @@ -0,0 +1,61 @@ +#!/usr/bin/env python3 +"""Isolate representative entry points without suppressing NVVM verification. + +Run each NVIDIA invocation in a child process so a verifier crash does not +prevent retaining evidence for the remaining kernels. No production IR changes. +""" +import argparse +import hashlib +import json +import os +from pathlib import Path +import re +import subprocess +import sys +from replay_cleanup import compile_nvvm + + +def main(): + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument('ir', type=Path) + parser.add_argument('--out', type=Path, required=True) + parser.add_argument('--worker', action='store_true') + args = parser.parse_args() + candidates = list(Path('target/cuda-builder-codegen').rglob('libintrinsics_v21.bc')) + unique = {hashlib.sha256(p.read_bytes()).hexdigest():p for p in candidates} + if len(unique) != 1: raise RuntimeError('expected one distinct LLVM 21 intrinsic library') + libraries = [Path(os.environ['CUDA_HOME'])/'nvvm/libdevice/libdevice.10.bc',next(iter(unique.values()))] + if args.worker: + compile_nvvm(args.ir,libraries,args.out) + return + args.out.mkdir(parents=True,exist_ok=True) + llvm_bin = Path(subprocess.check_output([os.environ['LLVM_CONFIG_21'],'--bindir'],text=True).strip()) + report = {'input_sha256':hashlib.sha256(args.ir.read_bytes()).hexdigest(), + 'libraries':{str(p):hashlib.sha256(p.read_bytes()).hexdigest() for p in libraries},'kernels':{}} + for kernel in ('wave2_float','wave2_shared','wave2_atomic','wave2_shuffle'): + out = args.out/kernel;out.mkdir(exist_ok=True) + globals_to_keep = set() + commands = [] + while True: + command = [str(llvm_bin/'llvm-extract'),'--recursive','--func='+kernel, + *['--glob='+s for s in sorted(globals_to_keep)],'-S',str(args.ir),'-o',str(out/'extracted.ll')] + commands.append(command);subprocess.run(command,check=True) + source = (out/'extracted.ll').read_text() + names = {s.strip('"') for s in re.findall(r'^@("[^"\\]*"|[-\w.$]+) = external (?:addrspace\(\d+\) )?',source,re.M)} + if not names: break + if names <= globals_to_keep: raise RuntimeError(f'unresolved globals: {names}') + globals_to_keep.update(names) + command = ['opt-21','-passes=verify',str(out/'extracted.ll'),'-o',str(out/'module.bc')] + commands.append(command);subprocess.run(command,check=True) + command = [sys.executable,str(Path(__file__).resolve()),str(out/'module.bc'),'--out',str(out),'--worker'] + commands.append(command) + with (out/'nvvm.log').open('w') as log: + result = subprocess.run(command,stdout=log,stderr=subprocess.STDOUT) + (out/'commands.json').write_text(json.dumps(commands,indent=2)+'\n') + report['kernels'][kernel] = {'exit_code':result.returncode,'compiled':result.returncode==0, + 'ir_sha256':hashlib.sha256((out/'extracted.ll').read_bytes()).hexdigest()} + (args.out/'diagnosis.json').write_text(json.dumps(report,indent=2)+'\n') + + +if __name__ == '__main__': + main() diff --git a/examples/ptx_export/representative-kernels/Cargo.toml b/examples/ptx_export/representative-kernels/Cargo.toml new file mode 100644 index 00000000..21b57e13 --- /dev/null +++ b/examples/ptx_export/representative-kernels/Cargo.toml @@ -0,0 +1,11 @@ +[package] +name = "ptx-representative-kernels" +version = "0.1.0" +edition = "2021" +publish = false + +[dependencies] +cuda_std = { path = "../../../crates/cuda_std" } + +[lib] +crate-type = ["cdylib", "rlib"] diff --git a/examples/ptx_export/representative-kernels/src/lib.rs b/examples/ptx_export/representative-kernels/src/lib.rs new file mode 100644 index 00000000..276dbc97 --- /dev/null +++ b/examples/ptx_export/representative-kernels/src/lib.rs @@ -0,0 +1,60 @@ +//! Runtime-input probes for floating point, shared memory, atomics and shuffles. +use core::{mem::MaybeUninit, sync::atomic::Ordering}; +use cuda_std::{address_space, kernel, thread, warp}; + +/// # Safety +/// Inputs and output address `count` elements; output does not overlap inputs. +#[kernel] +pub unsafe fn wave2_float(a: *const f32, b: *const f32, out: *mut f32, count: u32) { + let i = thread::index_1d(); + if i < count { + *out.add(i as usize) = *a.add(i as usize) * 0.5 + *b.add(i as usize); + } +} + +/// # Safety +/// Launch exactly 32 threads per block. Buffers address `count` u32s and do not overlap. +#[kernel] +pub unsafe fn wave2_shared(input: *const u32, out: *mut u32, count: u32) { + #[address_space(shared)] + static mut TILE: [MaybeUninit; 32] = [MaybeUninit::uninit(); 32]; + let i = thread::index_1d(); + let lane = thread::thread_idx_x() as usize; + let tile = core::ptr::addr_of_mut!(TILE).cast::(); + *tile.add(lane) = if i < count { *input.add(i as usize) } else { 0 }; + thread::sync_threads(); + if i < count { + *out.add(i as usize) = *tile.add(31 - lane); + } +} + +/// # Safety +/// `counter` addresses one aligned u32 initialized before launch. +#[kernel] +pub unsafe fn wave2_atomic(counter: *mut u32, count: u32) { + let i = thread::index_1d(); + if i < count { + cuda_std::atomic::mid::atomic_fetch_add_u32_device(counter, Ordering::Relaxed, i + 1); + } +} + +/// # Safety +/// Launch exactly 32 threads per block. Input addresses `count` u32s; output +/// addresses `4 * count` u32s. Buffers do not overlap. +/// All lanes, including tail padding, participate in the shuffle. +#[kernel] +pub unsafe fn wave2_shuffle(input: *const u32, out: *mut u32, count: u32) { + let i = thread::index_1d(); + let value = if i < count { *input.add(i as usize) } else { 0 }; + let results = [ + warp::warp_shuffle_idx(u32::MAX, value, 31, 32), + warp::warp_shuffle_up(u32::MAX, value, 1, 32), + warp::warp_shuffle_down(u32::MAX, value, 1, 32), + warp::warp_shuffle_xor(u32::MAX, value, 1, 32), + ]; + if i < count { + for (direction, (other, valid)) in results.into_iter().enumerate() { + *out.add(i as usize * 4 + direction) = if valid { other } else { u32::MAX }; + } + } +} diff --git a/examples/ptx_export/run_guarded_select.py b/examples/ptx_export/run_guarded_select.py new file mode 100644 index 00000000..a0f28ed2 --- /dev/null +++ b/examples/ptx_export/run_guarded_select.py @@ -0,0 +1,92 @@ +#!/usr/bin/env python3 +"""Numerically check runtime-input select kernels through a CUDA Driver API library. + +Native default: NVIDIA libcuda.so.1 and an unchanged PTX/cubin module. +An explicit compatible driver/module can test another consumer separately. +""" +import argparse +import ctypes as c +import os +from pathlib import Path + +MASK = (1 << 64) - 1 + + +def expected(table, limit, initial, filtered=False): + n = min(limit, 64) + odd_sum = sum(table[i] for i in range(1, n, 2)) & MASK + if filtered: + return [odd_sum, odd_sum] + observed = sum(table[(initial & 63) if i == 0 else ((i - 1) | 1)] for i in range(n)) & MASK + return [odd_sum, odd_sum, observed] + + +def main(): + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument('module', type=Path) + parser.add_argument('--kernel', choices=['rust_guarded_select', 'rust_filtered_select'], required=True) + parser.add_argument('--driver', default='libcuda.so.1') + args = parser.parse_args() + os.environ['CUMETAL_TRACE_GPU'] = '1' + os.environ['CUMETAL_ENABLE_WORKLOAD_SPECIALIZATIONS'] = '0' + lib = c.CDLL(args.driver) + ptr, u32, u64 = c.c_void_p, c.c_uint32, c.c_uint64 + + def api(name, types, *values): + # Prefer the size_t / 64-bit device-pointer versions on NVIDIA. + fn = getattr(lib, name + '_v2', None) or getattr(lib, name) + fn.argtypes, fn.restype = types, c.c_int + status = fn(*values) + if status: raise RuntimeError(f'{name} failed: {status}') + + api('cuInit', [u32], 0) + device = c.create_string_buffer(256) + api('cuDeviceGetName', [ptr, c.c_int, c.c_int], device, len(device), 0) + context, module, kernel = ptr(), ptr(), ptr() + api('cuCtxCreate', [c.POINTER(ptr), u32, c.c_int], c.byref(context), 0, 0) + total = 0 + try: + api('cuModuleLoad', [c.POINTER(ptr), c.c_char_p], c.byref(module), os.fsencode(args.module.resolve())) + api('cuModuleGetFunction', [c.POINTER(ptr), ptr, c.c_char_p], c.byref(kernel), module, args.kernel.encode()) + cases = [(limit, initial) for limit in [*range(66), (1 << 32) - 1] + for initial in [0, 1, 17, 63, 64, MASK]] + filtered = args.kernel == 'rust_filtered_select' + for seed in [0, 1, 1 << 63, MASK]: + table = [(i * 0x9e3779b97f4a7c15 + seed) & MASK for i in range(64)] + oracle = [v for limit, initial in cases for v in expected(table, limit, initial, filtered)] + poison = 0xa5a5a5a5a5a5a5a5 + result = (u64 * (len(oracle) + 16))(*([poison] * (len(oracle) + 16))) + inputs = [(u64 * 64)(*table), (u32 * len(cases))(*(x[0] for x in cases))] + if not filtered: inputs.append((u64 * len(cases))(*(x[1] for x in cases))) + allocations = [] + try: + for data in [*inputs, result]: + address = u64() + api('cuMemAlloc', [c.POINTER(u64), c.c_size_t], c.byref(address), c.sizeof(data)) + allocations.append(address) + api('cuMemcpyHtoD', [u64, ptr, c.c_size_t], address, c.cast(data, ptr), c.sizeof(data)) + # Native CUDA reads the low four bytes; wider backing storage + # also accommodates consumers whose scalar ABI reads eight. + count = u64(len(cases)) + parameters = (ptr * (len(allocations) + 2))( + *[c.cast(c.pointer(x), ptr) for x in [*allocations, count]], None) + api('cuLaunchKernel', [ptr] + [u32] * 7 + [ptr, c.POINTER(ptr), ptr], + kernel, (len(cases) + 63) // 64, 1, 1, 64, 1, 1, 0, None, parameters, None) + api('cuCtxSynchronize', []) + api('cuMemcpyDtoH', [ptr, u64, c.c_size_t], c.cast(result, ptr), allocations[-1], c.sizeof(result)) + for index, value in enumerate(oracle): + if result[index] != value: + raise RuntimeError(f'seed={seed} word={index}: got {result[index]:016x}, expected {value:016x}') + if list(result)[len(oracle):] != [poison] * 16: raise RuntimeError('output guard overwritten') + total += len(cases) + finally: + for address in reversed(allocations): api('cuMemFree', [u64], address) + print(f'NUMERICAL_PASS {args.kernel}: {total} cases, independent oracle and guards; ' + f'device={device.value.decode()} driver={args.driver}') + finally: + if module.value: api('cuModuleUnload', [ptr], module) + api('cuCtxDestroy', [ptr], context) + + +if __name__ == '__main__': + main() diff --git a/examples/ptx_export/run_wave2.py b/examples/ptx_export/run_wave2.py new file mode 100644 index 00000000..51ffa532 --- /dev/null +++ b/examples/ptx_export/run_wave2.py @@ -0,0 +1,102 @@ +#!/usr/bin/env python3 +"""Build one pinned workload and compare production DCE and bounded LLVM experiments.""" +import argparse +import hashlib +import json +from pathlib import Path +import re +import shutil +import subprocess +import sys +from replay_cleanup import normalized_functions + +MINER_COMMIT = '9791234249fc8cb762c296c4fda4503d2686ff77' +WORKLOADS = ('solana', 'bitcoin', 'ethereum', 'shallenge', 'self_test', 'representative', 'small') + + +def main(): + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument('workload', choices=WORKLOADS) + parser.add_argument('--out', type=Path, required=True) + parser.add_argument('--miner', type=Path) + args = parser.parse_args() + scripts = Path(__file__).resolve().parent + root = args.out.resolve(); root.mkdir(parents=True, exist_ok=True) + provenance = {'workload':args.workload, 'source_commit':subprocess.check_output(['git','rev-parse','HEAD'],text=True).strip()} + if args.workload in ('representative', 'small'): + kernels = scripts/('representative-kernels' if args.workload == 'representative' else 'kernels') + extra = [str(kernels)] + expected = set(re.findall(r'pub unsafe fn (\w+)\(', (kernels/'src/lib.rs').read_text())) + else: + if args.miner is None: raise RuntimeError('mining workload requires a pinned checkout') + miner = args.miner.resolve() + commit = subprocess.check_output(['git','-C',str(miner),'rev-parse','HEAD'],text=True).strip() + if commit != MINER_COMMIT: raise RuntimeError('mining revision differs from the experiment pin') + kernels = miner/'kernels' + manifest = kernels/'Cargo.toml'; original = manifest.read_text() + replacement = 'cuda_std = { path = '+json.dumps(str(scripts.parents[1]/'crates/cuda_std'))+' }' + patched, count = re.subn(r'^cuda_std = \{ git = "https://github.com/brandonros/Rust-CUDA.git", rev = "2f4fd1d" \}$',replacement,original,flags=re.M) + if count != 1: raise RuntimeError('unexpected cuda_std dependency') + (root/'Cargo.toml.original').write_text(original) + shutil.copy2(kernels/'Cargo.lock',root/'Cargo.lock.original') + manifest.write_text(patched); (root/'Cargo.toml.patched').write_text(patched) + provenance['mining_commit'] = commit + source_file = {'solana':'solana_vanity.rs','bitcoin':'bitcoin_vanity.rs','ethereum':'ethereum_vanity.rs','shallenge':'shallenge.rs','self_test':'self_test.rs'}[args.workload] + expected = set(re.findall(r'pub unsafe extern "C" fn (\w+)\(', (kernels/'src'/source_file).read_text())) + extra = [str(kernels),args.workload] + if not expected: raise RuntimeError('no expected kernel exports identified') + provenance['expected_entries'] = sorted(expected) + (root/'provenance.json').write_text(json.dumps(provenance,indent=2)+'\n') + modules = {} + failures = {} + for mode in ('none','default','inline-scalar'): + print(f'{args.workload}: building {mode}',flush=True) + out = root/mode; out.mkdir(parents=True,exist_ok=True) + command = ['cargo','run','-vv','-p','ptx_export','--features','llvm21','--',str(out),mode,*extra] + (out/'compiler-command.json').write_text(json.dumps(command,indent=2)+'\n') + with (out/'build.log').open('w') as log: + build = subprocess.run(command,stdout=log,stderr=subprocess.STDOUT) + if build.returncode: + failures[mode] = build.returncode + continue + subprocess.run(['opt-21','-passes=verify','-disable-output',str(out/'final-module.ll')],check=True) + print(f'{args.workload}: inspecting {mode}',flush=True) + subprocess.run([sys.executable,str(scripts/'inspect_codegen.py'),str(out), + '--reuse-from',str(root/'none')],check=True) + source = (out/'rust_kernels.ptx').read_text() + if args.workload == 'representative': + for direction in ('idx','up','down','bfly'): + if 'shfl.sync.'+direction not in source: + raise RuntimeError(f'{mode}: missing {direction} shuffle regression coverage') + entries = set(re.findall(r'\.entry\s+(\w+)\s*\(',source)) + if entries != expected: raise RuntimeError(f'{mode}: unexpected kernel exports: missing={expected-entries}, extra={entries-expected}') + modules[mode] = normalized_functions(source) + if failures: + (root/'build-failures.json').write_text(json.dumps(failures,indent=2)+'\n') + if args.workload == 'representative': + for mode in failures: + ir = root/mode/'final-module.ll' + if ir.exists(): + subprocess.run([sys.executable,str(scripts/'diagnose_wave2_nvvm.py'),str(ir), + '--out',str(root/mode/'isolated')],check=True) + raise RuntimeError(f'build failures (other modes were still checked): {failures}') + if modules['default'] != modules['none']: + raise RuntimeError('production default DCE changes baseline PTX function bodies; investigate before promotion') + if args.workload not in ('representative','small'): + shutil.copy2(kernels/'Cargo.lock',root/'Cargo.lock.resolved') + command = [sys.executable,str(scripts/'replay_cleanup.py'),str(root/'none'),'--wave2'] + if args.workload not in ('solana','small'): + command += ['--only','baseline,dce-only,inline-only'] + subprocess.run(command,check=True) + replay = root/'none/cleanup-experiment/inline-only/rust_kernels.ptx' + if normalized_functions(replay.read_text()) != modules['inline-scalar']: + raise RuntimeError('integrated InlineScalar differs from independent replay') + result = {'workload':args.workload,'expected_entry_count':len(expected),'default_matches_disabled_ptx':True, + 'inline_scalar_matches_replay':True,'nvidia_execution':False, + 'ptx_sha256':{m:hashlib.sha256((root/m/'rust_kernels.ptx').read_bytes()).hexdigest() for m in modules}} + (root/'comparison.json').write_text(json.dumps(result,indent=2)+'\n') + print(json.dumps(result),flush=True) + + +if __name__ == '__main__': + main() diff --git a/examples/ptx_export/run_wave2_gpu.py b/examples/ptx_export/run_wave2_gpu.py new file mode 100644 index 00000000..b5bb7545 --- /dev/null +++ b/examples/ptx_export/run_wave2_gpu.py @@ -0,0 +1,152 @@ +#!/usr/bin/env python3 +"""Numerical probes and opt-in CUDA-event timings for unchanged wave-2 modules. + +Use native NVIDIA libcuda.so.1 for timing. An explicit compatible driver can +validate another consumer numerically; those results are labeled separately. +""" +import argparse +import ctypes as c +import hashlib +import json +import os +from pathlib import Path +import statistics + +PTR, U32, U64 = c.c_void_p, c.c_uint32, c.c_uint64 +GUARD = 0xA5A5A5A5 + + +class Driver: + def __init__(self, library): + self.lib = c.CDLL(library) + + def call(self, name, types, *values): + fn = getattr(self.lib, name+'_v2', None) or getattr(self.lib, name) + fn.argtypes, fn.restype = types, c.c_int + status = fn(*values) + if status: + raise RuntimeError(f'{name}: CUDA status {status}') + + def launch(self, kernel, parameters, count): + self.call('cuLaunchKernel', [PTR]+[U32]*7+[PTR,c.POINTER(PTR),PTR], + kernel, (count+31)//32,1,1,32,1,1,0,None,parameters,None) + + +def input_and_oracle(kernel, count): + if kernel == 'wave2_float': + # Binary fractions in this range are exact with either fused or separate operations. + a = [(i % 1024 - 512) * 0.25 for i in range(count)] + b = [(i % 257 - 128) * 0.5 for i in range(count)] + expected = [(c.cast(c.pointer(c.c_float(x*0.5+y)), c.POINTER(U32))[0]) for x,y in zip(a,b)] + inputs = [(c.c_float*count)(*a), (c.c_float*count)(*b)] + elif kernel == 'wave2_atomic': + inputs, expected = [], [(count*(count+1)//2) & 0xFFFFFFFF] + else: + data = [(i*2654435761+17) & 0xFFFFFFFF for i in range(count)] + if kernel == 'wave2_shared': + partners = [(i//32)*32+31-i%32 for i in range(count)] + expected = [data[p] if p < count else 0 for p in partners] + else: + expected = [] + for i in range(count): + base, lane = i//32*32, i%32 + # IDX(31), UP(1), DOWN(1), XOR(1), including invalid edge lanes. + for partner in (31,lane-1,lane+1,lane^1): + expected.append(0xFFFFFFFF if not 0<=partner<32 else + data[base+partner] if base+partner2**24 for n in counts) or not 1<=args.repeats<=10000: + parser.error('counts must be 1..2^24 and repeats 1..10000') + os.environ['CUMETAL_ENABLE_WORKLOAD_SPECIALIZATIONS']='0' + os.environ['CUMETAL_TRACE_GPU']='1' + driver=Driver(args.driver) + driver.call('cuInit',[U32],0) + device=c.create_string_buffer(256) + driver.call('cuDeviceGetName',[PTR,c.c_int,c.c_int],device,len(device),0) + version=c.c_int();driver.call('cuDriverGetVersion',[c.POINTER(c.c_int)],c.byref(version)) + context,module,kernel=PTR(),PTR(),PTR() + driver.call('cuCtxCreate',[c.POINTER(PTR),U32,c.c_int],c.byref(context),0,0) + report={'module_sha256':hashlib.sha256(args.module.read_bytes()).hexdigest(),'runner_sha256':hashlib.sha256(Path(__file__).read_bytes()).hexdigest(), + 'driver':args.driver,'driver_version':version.value,'device':device.value.decode(), + 'native_nvidia_driver_requested':args.driver=='libcuda.so.1','results':[]} + args.out.parent.mkdir(parents=True,exist_ok=True) + try: + driver.call('cuModuleLoad',[c.POINTER(PTR),c.c_char_p],c.byref(module),os.fsencode(args.module.resolve())) + driver.call('cuModuleGetFunction',[c.POINTER(PTR),PTR,c.c_char_p],c.byref(kernel),module,args.kernel.encode()) + for count in counts: + result=run_case(driver,kernel,args.kernel,count,args.repeats if args.benchmark else 0) + report['results'].append(result) + args.out.write_text(json.dumps(report,indent=2)+'\n') + print(json.dumps(result),flush=True) + except Exception as error: + report['error']=str(error);args.out.write_text(json.dumps(report,indent=2)+'\n');raise + finally: + if module.value: driver.call('cuModuleUnload',[PTR],module) + driver.call('cuCtxDestroy',[PTR],context) + + +if __name__=='__main__': + main() diff --git a/examples/ptx_export/test_module_cleanup.py b/examples/ptx_export/test_module_cleanup.py new file mode 100644 index 00000000..45a26114 --- /dev/null +++ b/examples/ptx_export/test_module_cleanup.py @@ -0,0 +1,28 @@ +import unittest +from check_module_cleanup import canonical + + +class ModuleComparisonTests(unittest.TestCase): + def test_comments_do_not_change_ir(self): + self.assertEqual(canonical('; ModuleID = one\nbb: ; preds = %a, %b\n ret void\n'), + canonical('; ModuleID = two\nbb: ; preds = %b, %a\n ret void\n')) + + def test_strings_and_semantic_differences_are_preserved(self): + line = '@"semi;colon" = constant [2 x i8] c";x" ; comment' + self.assertEqual(canonical(line), '@"semi;colon" = constant [2 x i8] c";x"') + self.assertNotEqual(canonical('ret i32 1 ; first'), canonical('ret i32 2 ; second')) + self.assertNotEqual(canonical('@s = constant [1 x i8] c";"'), + canonical('@s = constant [1 x i8] c":"')) + + def test_phi_order_preserves_value_block_associations(self): + a = '%0 = phi i64 [ 3, %1 ], [ 5, %2 ]' + b = '%0 = phi i64 [ 5, %2 ], [ 3, %1 ]' + wrong = '%0 = phi i64 [ 5, %1 ], [ 3, %2 ]' + self.assertEqual(canonical(a), canonical(b)) + self.assertNotEqual(canonical(a), canonical(wrong)) + self.assertEqual(canonical('%0 = phi [2 x i8] [ [i8 1, i8 2], %1 ], [ zeroinitializer, %2 ]'), + canonical('%0 = phi [2 x i8] [ zeroinitializer, %2 ], [ [i8 1, i8 2], %1 ]')) + + +if __name__ == '__main__': + unittest.main() From f22a716e40aaa9ea45732174e15b9bc4cfd0352c Mon Sep 17 00:00:00 2001 From: Brandon Ros Date: Mon, 14 Sep 2026 21:12:40 -0400 Subject: [PATCH 2/2] Isolate pinned mining checkouts and clean up after failed experiments --- examples/ptx_export/check_workloads.py | 27 ++++------ examples/ptx_export/run_wave2.py | 24 ++++----- examples/ptx_export/test_workload.py | 75 ++++++++++++++++++++++++++ examples/ptx_export/workload.py | 41 ++++++++++++++ 4 files changed, 137 insertions(+), 30 deletions(-) create mode 100644 examples/ptx_export/test_workload.py create mode 100644 examples/ptx_export/workload.py diff --git a/examples/ptx_export/check_workloads.py b/examples/ptx_export/check_workloads.py index ec77874e..5b21c7f0 100644 --- a/examples/ptx_export/check_workloads.py +++ b/examples/ptx_export/check_workloads.py @@ -1,13 +1,14 @@ #!/usr/bin/env python3 """Build size-oriented reproducers or pinned, unmodified Solana kernel sources.""" import argparse +from contextlib import nullcontext import json from pathlib import Path -import re import shutil import subprocess import sys from replay_cleanup import normalized_functions +from workload import MINER_COMMIT, prepared_miner def main(): @@ -17,23 +18,15 @@ def main(): args = parser.parse_args() root = args.artifacts.resolve() scripts = Path(__file__).resolve().parent + context = prepared_miner(args.miner, root/'mining-solana') if args.miner else nullcontext(None) + with context as miner: + compare(args, root, scripts, miner) + + +def compare(args, root, scripts, miner): if args.miner: - miner = args.miner.resolve() - out = root/'mining-solana'; out.mkdir(parents=True, exist_ok=True) - commit = subprocess.check_output(['git', '-C', str(miner), 'rev-parse', 'HEAD'], text=True).strip() - if commit != '9791234249fc8cb762c296c4fda4503d2686ff77': - raise RuntimeError('unexpected mining workload revision') - manifest = miner/'kernels/Cargo.toml' - original = manifest.read_text() - shutil.copy2(manifest, out/'Cargo.toml.original') - shutil.copy2(miner/'kernels/Cargo.lock', out/'Cargo.lock.original') - # Use this backend's cuda_std while retaining the mining algorithm source. - replacement = 'cuda_std = { path = '+json.dumps(str(scripts.parents[1]/'crates/cuda_std'))+' }' - patched, count = re.subn(r'^cuda_std = \{ git = "https://github.com/brandonros/Rust-CUDA.git", rev = "2f4fd1d" \}$', replacement, original, flags=re.M) - if count != 1: raise RuntimeError('unexpected cuda_std dependency; review workload pin') - manifest.write_text(patched) - shutil.copy2(manifest, out/'Cargo.toml.patched') - (out/'workload.json').write_text(json.dumps({'repository':'brandonros/vanity-miner-rs', 'commit':commit, + out = root/'mining-solana' + (out/'workload.json').write_text(json.dumps({'repository':'brandonros/vanity-miner-rs', 'commit':MINER_COMMIT, 'features':['solana'], 'source_change':'cuda_std dependency path only; algorithm source unchanged'}, indent=2)+'\n') builds = [('none', out, [str(miner/'kernels'), 'solana'])] else: diff --git a/examples/ptx_export/run_wave2.py b/examples/ptx_export/run_wave2.py index 51ffa532..28f7e360 100644 --- a/examples/ptx_export/run_wave2.py +++ b/examples/ptx_export/run_wave2.py @@ -1,6 +1,7 @@ #!/usr/bin/env python3 """Build one pinned workload and compare production DCE and bounded LLVM experiments.""" import argparse +from contextlib import nullcontext import hashlib import json from pathlib import Path @@ -10,7 +11,7 @@ import sys from replay_cleanup import normalized_functions -MINER_COMMIT = '9791234249fc8cb762c296c4fda4503d2686ff77' +from workload import MINER_COMMIT, prepared_miner WORKLOADS = ('solana', 'bitcoin', 'ethereum', 'shallenge', 'self_test', 'representative', 'small') @@ -22,25 +23,22 @@ def main(): args = parser.parse_args() scripts = Path(__file__).resolve().parent root = args.out.resolve(); root.mkdir(parents=True, exist_ok=True) + if args.workload not in ('representative', 'small') and args.miner is None: + parser.error('mining workload requires a pinned checkout') + context = prepared_miner(args.miner, root) if args.workload not in ('representative', 'small') else nullcontext(None) + with context as miner: + compare(args, scripts, root, miner) + + +def compare(args, scripts, root, miner): provenance = {'workload':args.workload, 'source_commit':subprocess.check_output(['git','rev-parse','HEAD'],text=True).strip()} if args.workload in ('representative', 'small'): kernels = scripts/('representative-kernels' if args.workload == 'representative' else 'kernels') extra = [str(kernels)] expected = set(re.findall(r'pub unsafe fn (\w+)\(', (kernels/'src/lib.rs').read_text())) else: - if args.miner is None: raise RuntimeError('mining workload requires a pinned checkout') - miner = args.miner.resolve() - commit = subprocess.check_output(['git','-C',str(miner),'rev-parse','HEAD'],text=True).strip() - if commit != MINER_COMMIT: raise RuntimeError('mining revision differs from the experiment pin') kernels = miner/'kernels' - manifest = kernels/'Cargo.toml'; original = manifest.read_text() - replacement = 'cuda_std = { path = '+json.dumps(str(scripts.parents[1]/'crates/cuda_std'))+' }' - patched, count = re.subn(r'^cuda_std = \{ git = "https://github.com/brandonros/Rust-CUDA.git", rev = "2f4fd1d" \}$',replacement,original,flags=re.M) - if count != 1: raise RuntimeError('unexpected cuda_std dependency') - (root/'Cargo.toml.original').write_text(original) - shutil.copy2(kernels/'Cargo.lock',root/'Cargo.lock.original') - manifest.write_text(patched); (root/'Cargo.toml.patched').write_text(patched) - provenance['mining_commit'] = commit + provenance['mining_commit'] = MINER_COMMIT source_file = {'solana':'solana_vanity.rs','bitcoin':'bitcoin_vanity.rs','ethereum':'ethereum_vanity.rs','shallenge':'shallenge.rs','self_test':'self_test.rs'}[args.workload] expected = set(re.findall(r'pub unsafe extern "C" fn (\w+)\(', (kernels/'src'/source_file).read_text())) extra = [str(kernels),args.workload] diff --git a/examples/ptx_export/test_workload.py b/examples/ptx_export/test_workload.py new file mode 100644 index 00000000..19bedc3f --- /dev/null +++ b/examples/ptx_export/test_workload.py @@ -0,0 +1,75 @@ +from pathlib import Path +import subprocess +import sys +import tempfile +import unittest +from unittest.mock import patch + +import check_workloads +import run_wave2 +import workload + + +class WorkloadTests(unittest.TestCase): + def setUp(self): + self.directory = tempfile.TemporaryDirectory() + self.addCleanup(self.directory.cleanup) + self.root = Path(self.directory.name) + self.source = self.root / 'source' + kernels = self.source / 'kernels' + kernels.mkdir(parents=True) + self.manifest = kernels / 'Cargo.toml' + self.original = 'cuda_std = { git = "https://github.com/brandonros/Rust-CUDA.git", rev = "2f4fd1d" }\n' + self.manifest.write_text(self.original) + (kernels / 'Cargo.lock').write_text('committed lock\n') + self.git('init', '--quiet') + self.git('add', '.') + self.git('-c', 'user.name=Test', '-c', 'user.email=test@example.invalid', + '-c', 'commit.gpgsign=false', 'commit', '--quiet', '-m', 'fixture') + self.pin = self.git('rev-parse', 'HEAD').strip() + self.manifest.write_text('local edits\n') + (kernels / 'Cargo.lock').write_text('local lock edits\n') + (kernels / 'untracked').write_text('keep me\n') + + def git(self, *args): + return subprocess.check_output(['git', '-C', str(self.source), *args], text=True) + + def assert_source_unchanged(self): + self.assertEqual(self.manifest.read_text(), 'local edits\n') + self.assertEqual((self.manifest.parent / 'Cargo.lock').read_text(), 'local lock edits\n') + self.assertEqual((self.manifest.parent / 'untracked').read_text(), 'keep me\n') + + def test_repeated_preparation_uses_committed_source_and_cleans_up(self): + with patch.object(workload, 'MINER_COMMIT', self.pin): + for _ in range(2): + with workload.prepared_miner(self.source, self.root / 'artifacts') as miner: + self.assertIn('path = ', (miner / 'kernels/Cargo.toml').read_text()) + self.assertFalse((miner / 'kernels/untracked').exists()) + (miner / 'kernels/Cargo.lock').write_text('resolved lock\n') + self.assert_source_unchanged() + self.assertFalse(miner.exists()) + self.assertEqual((self.root / 'artifacts/Cargo.toml.original').read_text(), self.original) + self.assert_source_unchanged() + + def test_both_runners_cleanup_after_failure_and_allow_retry(self): + def fail(*args): + self.prepared = args[-1] + self.assertTrue(self.prepared.exists()) + raise RuntimeError('build failed') + + for module in (run_wave2, check_workloads): + args = ['runner', 'solana', '--out', str(self.root / 'out')] if module is run_wave2 else ['runner', str(self.root / 'out')] + args += ['--miner', str(self.source)] + with patch.object(workload, 'MINER_COMMIT', self.pin), patch.object(sys, 'argv', args), patch.object(module, 'compare', side_effect=fail): + for _ in range(2): + with self.assertRaisesRegex(RuntimeError, 'build failed'): + module.main() + self.assertFalse(self.prepared.exists()) + self.assert_source_unchanged() + + def test_wrong_revision_is_rejected_without_modifying_source(self): + with patch.object(workload, 'MINER_COMMIT', '0' * 40): + with self.assertRaisesRegex(RuntimeError, 'revision differs'): + with workload.prepared_miner(self.source, self.root / 'artifacts'): + self.fail('accepted incorrect revision') + self.assert_source_unchanged() diff --git a/examples/ptx_export/workload.py b/examples/ptx_export/workload.py new file mode 100644 index 00000000..62cf519a --- /dev/null +++ b/examples/ptx_export/workload.py @@ -0,0 +1,41 @@ +"""Prepare the pinned mining source without modifying the supplied checkout.""" +from contextlib import contextmanager +import json +from pathlib import Path +import re +import shutil +import subprocess +import tempfile + +MINER_COMMIT = '9791234249fc8cb762c296c4fda4503d2686ff77' + + +@contextmanager +def prepared_miner(source, artifacts): + source = source.resolve() + commit = subprocess.check_output( + ['git', '-C', str(source), 'rev-parse', 'HEAD'], text=True).strip() + if commit != MINER_COMMIT: + raise RuntimeError('mining revision differs from the experiment pin') + artifacts.mkdir(parents=True, exist_ok=True) + with tempfile.TemporaryDirectory(prefix='ptx-miner-') as directory: + miner = Path(directory) / 'miner' + # Read committed sources only; local edits and untracked files are not inputs. + subprocess.run(['git', 'clone', '--quiet', '--shared', '--no-checkout', + str(source), str(miner)], check=True) + subprocess.run(['git', '-C', str(miner), 'checkout', '--quiet', '--detach', + commit], check=True) + manifest = miner / 'kernels/Cargo.toml' + original = manifest.read_text() + cuda_std = Path(__file__).resolve().parents[2] / 'crates/cuda_std' + replacement = 'cuda_std = { path = ' + json.dumps(str(cuda_std)) + ' }' + patched, count = re.subn( + r'^cuda_std = \{ git = "https://github.com/brandonros/Rust-CUDA.git", rev = "2f4fd1d" \}$', + lambda _: replacement, original, flags=re.M) + if count != 1: + raise RuntimeError('unexpected cuda_std dependency; review workload pin') + (artifacts / 'Cargo.toml.original').write_text(original) + shutil.copy2(miner / 'kernels/Cargo.lock', artifacts / 'Cargo.lock.original') + manifest.write_text(patched) + (artifacts / 'Cargo.toml.patched').write_text(patched) + yield miner