From 221e1f13f758ebe8d9f188bab976c361f7cacf45 Mon Sep 17 00:00:00 2001 From: Digital Drosophila Archive Date: Thu, 1 Oct 2026 14:41:59 +0800 Subject: [PATCH] Preserve exact float32 via explicit nvcc compilation and audited binary loading --- .../conditioning-acceleration-v1/PROTOCOL.md | 6 + .../finalize_report.py | 109 ++++++++++++++ .../gpu_backend.py | 140 ++++++++++++++++-- .../test_gpu_math.py | 84 ++++++++++- 4 files changed, 326 insertions(+), 13 deletions(-) create mode 100644 models/male-cns-lif/conditioning-acceleration-v1/finalize_report.py diff --git a/models/male-cns-lif/conditioning-acceleration-v1/PROTOCOL.md b/models/male-cns-lif/conditioning-acceleration-v1/PROTOCOL.md index 1882606..1359ea3 100644 --- a/models/male-cns-lif/conditioning-acceleration-v1/PROTOCOL.md +++ b/models/male-cns-lif/conditioning-acceleration-v1/PROTOCOL.md @@ -23,3 +23,9 @@ Budget: at most 15 minutes of candidate benchmark execution after dependencies/c Passing this benchmark means only ELIGIBILITY-UPDATE-REPLAY-EQUIVALENCE for the tested inputs. It does not establish full-brain event equivalence, end-to-end acceleration, learning success, resumability, or completion of Phase 4AR. Integration, persistent brain workers and independent-seed multi-GPU scheduling require separately validated execution/audit changes. Stop before installing any update or starting learning. Keep source, configuration, environment installation report, fixed input manifest, raw timing samples, comparisons, failures and SHA-256 manifest. Check Phase 4AR registration/source and the prior saved-file inventory before/after; report the exact scope (not a new Phase 1–4AQ numerical regression). + +## Recorded implementation correction (before successful GPU execution) + +The first candidate installation attempted pip's `--report` with the new venv's inherited pip 22.0.2, which does not support it. Only the candidate venv was upgraded to pip 25.2; both logs remain preserved. + +GPU attempt `v100a-01` failed before running kernels because CuPy 13.6's RawKernel cache compiler appends `-ftz=true`, conflicting with the required `--ftz=false`. Do not remove the precision requirement. The corrected candidate compiles the same CUDA recurrence with installed nvcc 12.6, explicit separate rounding and `--ftz=false`, then loads the precompiled module through CuPy. Compiler commands, sources, output and binary hashes are preserved in candidate-only caches. The original scientific recurrence and all predeclared fixtures remain unchanged; failed attempt records are not replaced. diff --git a/models/male-cns-lif/conditioning-acceleration-v1/finalize_report.py b/models/male-cns-lif/conditioning-acceleration-v1/finalize_report.py new file mode 100644 index 0000000..7136453 --- /dev/null +++ b/models/male-cns-lif/conditioning-acceleration-v1/finalize_report.py @@ -0,0 +1,109 @@ +"""Aggregate measured replay-only results; never estimate unmeasured full-training acceleration.""" +import argparse +from datetime import datetime, timezone +import hashlib +import json +from pathlib import Path +import statistics +import subprocess + +BASE = Path('/home/qishizhe/phase4ar-v100-acceleration-20261001') +SOURCE = Path(__file__).resolve().parent + +def read(path): return json.loads(Path(path).read_text()) +def sha(path): + h = hashlib.sha256() + with Path(path).open('rb') as f: + for c in iter(lambda: f.read(8 << 20), b''): h.update(c) + return h.hexdigest() +def save(path, data): + with Path(path).open('x') as f: json.dump(data, f, indent=2, sort_keys=True, allow_nan=False) + +p = argparse.ArgumentParser() +p.add_argument('--run', action='append', required=True) +a = p.parse_args() +assert len(a.run) == len(set(a.run)) == 3 +integrity = read(BASE / 'integrity-after-gpu.json') +assert integrity['status'] == 'PASS' +devices = [] +pins = set() +for name in a.run: + path = BASE / name + assert path.parent == BASE and path.is_dir() + s = read(path / 'summary.json') + assert s['status'] == 'PASS-ELIGIBILITY-UPDATE-REPLAY-EQUIVALENCE' + assert s['fixtures'] == 30 and s['rejected_invalid_cases'] == 6 and s['all_fixture_repeats_bitwise_exact'] + process = read(BASE / 'jobs' / (name + '.exit.json')) + assert process['observed_exit'] and process['exit_code'] == 0 and not process['timed_out'] + pins.add(s['candidate_commit']) + real = [read(f) for f in sorted(path.glob('fixtures--*.json'))] + assert len(real) == 14 + medians = [r['summaries']['replay_wall_including_preparation_and_comparison_seconds']['median'] for r in real] + kernels = [statistics.median(t['kernel_seconds'] for t in r['resident_timing']['resident_kernel_samples']) for r in real] + cpu = [r['reference_timing']['warm_update_seconds'] for r in real] + cpu_cold = [r['reference_timing']['cold_update_seconds'] for r in real] + devices.append(dict(run=name, device=s['device'], real_fixture_count=len(real), + median_real_fixture_cpu_original_update_api_seconds=statistics.median(cpu), + median_real_fixture_cpu_cold_update_api_seconds=statistics.median(cpu_cold), + median_real_fixture_candidate_replay_seconds=statistics.median(medians), + median_real_fixture_resident_kernel_seconds=statistics.median(kernels), + mean_real_fixture_candidate_replay_seconds=statistics.mean(medians), + original_update_api_to_candidate_replay_ratio=statistics.median([x/y for x,y in zip(cpu,medians)]), + ratio_warning='Different API overhead boundaries: original U checks/AST every call; candidate authenticates sources/mapping once. NOT pure CPU/GPU hardware ratio.', + cpu_baseline_sample_warning='One cold plus ONE warm original API sample per fixture; GPU 3 warmups +20 repeats per fixture.', + first_use_compile_seconds=real[0]['first_call_timing']['cold_compile_seconds'], + context_initialization_seconds=s['device']['initialization_seconds'], + observed_cuda_pool_reserved_peak_bytes=s['observed_cuda_pool_reserved_peak_bytes'], + positive_fixtures=30, invalid_inputs_rejected=6, bitwise_exact=True, + benchmark_wall_seconds=s['wall_seconds'], result_sha256=sha(path / 'summary.json'))) +assert len(pins) == 1, 'GPU runs used different candidate commits' +summary = dict(status='PASS-ELIGIBILITY-UPDATE-REPLAY-EQUIVALENCE', + created_utc=datetime.now(timezone.utc).isoformat(), benchmark_candidate_commit=next(iter(pins)), + report_code_commit=subprocess.check_output(['git','-C',str(SOURCE),'rev-parse','HEAD'],text=True).strip(), + devices=devices, frozen_integrity=integrity, + scientific_experiment_status='STOPPED, unchanged:1033/30000 CPU-verified training epochs and138/3920 probes', + original_experiment_resumed=False, full_training_acceleration_measured=False, weights_installed=False, + full_brain_gpu_event_equivalence_tested=False, learned_candidate_created=False, + limitations=['Replay module only, not a production training backend', + 'The sampled committed prefix covers7arms but only available physiological/PAM15/reinforce_A subset', + 'Repeated kernel calculations reset state and replay identical saved inputs, not additional learning trials', + 'Original CPU verification remains independent and required', + 'No simultaneous two-V100 training/scheduling benchmark; devices tested separately', + 'Original9.07-day estimate is not replaced by this component benchmark'], + fixes_preserved=['Initial candidate pip22 lacked--report; upgraded pip only in new venv', + 'First GPU compile rejected conflictingFTZ options; no kernel executed; explicit nvcc compilation retains original precision'], + next_work='Separately validate persistent training/update/verification workers and complete chain timing before authorizing a new scientific execution version') +save(BASE / 'final-summary.json', summary) +lines = ['# V100 conditioning acceleration — engineering replay result', '', + 'Status: **PASS-ELIGIBILITY-UPDATE-REPLAY-EQUIVALENCE**. This is not a learning result or a production-training promotion.', '', + f'Benchmark code commit: `{next(iter(pins))}`.', '', + '| Device | Real-fixture replay median (prep + transfers + comparison) | Resident CUDA kernel median |', + '|---|---:|---:|'] +for d in devices: + lines.append(f"| {d['device']['runtime_name']} ({d['run']}) | {d['median_real_fixture_candidate_replay_seconds']*1000:.3f} ms | {d['median_real_fixture_resident_kernel_seconds']*1000:.3f} ms |") +lines += ['', 'Each GPU passed14sealed real fixtures +16fixed synthetic cases, with sampled trace, final trace, allUS snapshots and delta byte-identical to the original independent CPU oracle. Six invalid inputs were rejected on each GPU.', '', + 'CPU warm reference includes the original API\'s repeated provenance checks and AST parsing; it is only one warm sample perfixture. GPU authenticates the source and complete anatomical mapping once. Therefore ratios in JSON describe API implementation replay, not pure hardware acceleration.', '', + 'The CUDA kernels preserve every timestep, float32 separate round-to-nearest operations, exact US timing, noUS zeros, compartment/frozen/D0 identity and the depression floor. No gain/learning-rate/tau/threshold/connectome/sign changes. No dense time×neuron matrix.', '', + 'A new candidate-only CuPy environment and explicit CUDA12.6 nvcc compilation are used. Original science and ComfyUI environments are unchanged. All three ComfyUI queues were checked before each fixture; no services were stopped.', '', + f"Integrity: {integrity['files_checked']} saved files, {integrity['bytes_checked']} bytes, {integrity['sources_checked']} registered source files checked; no mismatch. This is the saved Phase4AR inventory, not a new Phase1–4AQ numerical regression.", '', + 'Phase4AR remains stopped at1033/30000CPU-verified epochs and138/3920probes. No weights were installed and no training was resumed.', '', + 'Next: a separately validated persistent-worker pipeline and end-to-end benchmark. This local optimization does not justify dividing the9-day estimate by the kernel speedup. Independent-seed V100 parallelism is not yet measured.', '', + 'Preserved failures: candidate packaging incompatibility (--report unsupported bypip22); first CuPy compile FTZ flag conflict (before any kernel ran). Both were corrected only in the independent branch.', '', + 'Compiler references: [CuPy precompiled RawModule](https://docs.cupy.dev/en/v13.6.0/reference/generated/cupy.RawModule.html), [CUDA12.6 compiler options](https://docs.nvidia.com/cuda/archive/12.6.3/cuda-compiler-driver-nvcc/index.html).', ''] +with (BASE / 'RESULTS.md').open('x') as f: f.write('\n'.join(lines)) +# Exclude installed environment, reproducible compiler caches and final manifests themselves; +# environment packages/wheel hashes are retained in pip-install-report + freeze. +excluded = {'.venv'} +files = {} +for path in sorted(BASE.rglob('*')): + rel = path.relative_to(BASE) + if path.is_file() and not path.is_symlink() and not any(part in excluded for part in rel.parts): + files[rel.as_posix()] = dict(sha256=sha(path), bytes=path.stat().st_size) +save(BASE / 'artifact-manifest.json', dict(schema='independent-acceleration-artifacts-v1', + code_commit=summary['report_code_commit'], source_files={p.name:sha(p) for p in SOURCE.iterdir() if p.is_file()}, + files=files, excluded=['.venv (package versions and downloaded wheel hashes are in installation receipts)'], + no_frozen_source_modified=True, no_training_run=True)) +with (BASE / 'SHA256SUMS').open('x') as f: + for rel, info in files.items(): f.write(info['sha256'] + ' ' + rel + '\n') + f.write(sha(BASE/'artifact-manifest.json') + ' artifact-manifest.json\n') +print(json.dumps(summary), flush=True) diff --git a/models/male-cns-lif/conditioning-acceleration-v1/gpu_backend.py b/models/male-cns-lif/conditioning-acceleration-v1/gpu_backend.py index dec94db..867614e 100644 --- a/models/male-cns-lif/conditioning-acceleration-v1/gpu_backend.py +++ b/models/male-cns-lif/conditioning-acceleration-v1/gpu_backend.py @@ -3,7 +3,8 @@ The caller authenticates frozen sources, configuration, the complete KC list, and the candidate edge mapping. This module does not infer D from DAN activity. CuPy is imported only when CudaEligibility is instantiated. Set CUPY_CACHE_DIR -to a dedicated candidate directory before doing so; this module writes no files. +to a dedicated candidate directory before doing so. The only files this module +writes are explicit nvcc build evidence and cubins beneath that cache directory. API: prepare_events(...) once, then CudaEligibility(...).run(prepared, ...). All commands are sampled AFTER the current tick; only the LAST command snapshot @@ -14,9 +15,12 @@ replacement for comparison against the SHA-pinned frozen implementation. from __future__ import annotations from dataclasses import dataclass +import hashlib +import json import math import os from pathlib import Path +import subprocess import time from typing import Any @@ -26,6 +30,7 @@ import numpy as np COMPARTMENTS = ("B2B'2a", "a1", "y2a'1", "y5B'2a") _INT32_MAX = np.iinfo(np.int32).max CUDA_OPTIONS = ("--std=c++11", "--fmad=false", "--ftz=false") +DEFAULT_NVCC = "/usr/local/cuda/bin/nvcc" CUDA_SOURCE = r''' extern "C" __global__ void eligibility_trace(const long long* ptr, const int* ticks, @@ -84,6 +89,105 @@ void eligibility_delta(const float* sampled, const int* source_local, ''' +def _file_sha256(path: Path) -> str: + digest = hashlib.sha256() + with path.open("rb") as stream: + for block in iter(lambda: stream.read(1024 * 1024), b""): + digest.update(block) + return digest.hexdigest() + + +def _save_json_new(path: Path, value: Any) -> None: + with path.open("x", encoding="utf-8", newline="\n") as stream: + json.dump(value, stream, sort_keys=True, indent=2, allow_nan=False) + stream.write("\n") + + +def _nvcc_cubin(cache_root: Path, capability: str, nvcc_path: Path) -> tuple[Path, dict, dict]: + """Compile outside CuPy's compiler path, which injects -ftz=true in 13.6. + + A cache entry is immutable. Failed/incomplete entries are retained and cause + an error on reuse; the caller can select a fresh candidate cache directory. + Only the subprocess's temporary-directory environment is changed. + """ + probe_started = time.perf_counter() + _require(capability.isdigit() and 2 <= len(capability) <= 3, "invalid CUDA compute capability") + _require(cache_root.is_absolute(), "candidate cache directory must be absolute") + _require(nvcc_path.is_absolute() and nvcc_path.is_file(), "nvcc must be an existing absolute compiler path") + _require(not any(os.environ.get(key, "").strip() + for key in ("NVCC_PREPEND_FLAGS", "NVCC_APPEND_FLAGS")), + "NVCC injected flags are incompatible with explicit candidate compiler options") + cache_root = cache_root.resolve() + nvcc_path = nvcc_path.resolve() + version = subprocess.run([str(nvcc_path), "--version"], capture_output=True, text=True, + timeout=30, check=False) + _require(version.returncode == 0, "nvcc --version failed: " + version.stderr) + options = ("--cubin", "-arch=sm_" + capability, *CUDA_OPTIONS, "--keep") + identity = dict(schema="conditioning-cuda-nvcc-build-v1", + source_sha256=hashlib.sha256(CUDA_SOURCE.encode("utf-8")).hexdigest(), + nvcc_path=str(nvcc_path), nvcc_sha256=_file_sha256(nvcc_path), + nvcc_version_stdout=version.stdout, nvcc_version_stderr=version.stderr, + options=list(options), compute_capability=capability, + nvcc_ccbin=os.environ.get("NVCC_CCBIN", "")) + key = hashlib.sha256(json.dumps(identity, sort_keys=True, separators=(",", ":")).encode()).hexdigest() + build_dir = cache_root / "eligibility-nvcc" / key + _require(not build_dir.is_symlink() and build_dir.resolve().is_relative_to(cache_root), + "candidate build cache path escapes the requested directory") + cubin = build_dir / "eligibility.cubin" + receipt_path = build_dir / "build-receipt.json" + probe_seconds = time.perf_counter() - probe_started + if build_dir.exists(): + _require(receipt_path.is_file(), "incomplete nvcc candidate build retained: " + str(build_dir)) + receipt = json.loads(receipt_path.read_text(encoding="utf-8")) + _require(receipt.get("complete") is True and receipt["identity"] == identity, + "failed or mismatched nvcc candidate build retained: " + str(build_dir)) + for name, expected in receipt["files_sha256"].items(): + path = build_dir / name + _require(path.name == name and not path.is_symlink() + and path.resolve().is_relative_to(build_dir) + and _file_sha256(path) == expected, "nvcc cache artifact changed: " + name) + _require(receipt["files_sha256"].get("eligibility.cubin") == _file_sha256(cubin) + and receipt["files_sha256"].get("eligibility.cu") == identity["source_sha256"], + "nvcc cache lacks authenticated source/cubin") + return cubin, dict(cache_key=key, cache_directory=str(build_dir), cache_hit=True, + receipt_sha256=_file_sha256(receipt_path), **receipt), dict( + compiler_probe_seconds=probe_seconds, nvcc_compile_seconds=0.0) + build_dir.mkdir(parents=True, exist_ok=False) + source_path = build_dir / "eligibility.cu" + with source_path.open("x", encoding="utf-8", newline="\n") as stream: + stream.write(CUDA_SOURCE) + command = [str(nvcc_path), *options, "--keep-dir", str(build_dir), + str(source_path), "-o", str(cubin)] + _save_json_new(build_dir / "compile-command.json", dict(argv=command, cwd=str(build_dir), + temp_directory=str(build_dir), inherited_nvcc_injected_flags=False)) + _save_json_new(build_dir / "compiler-identity.json", identity) + environment = dict(os.environ, TMPDIR=str(build_dir), TMP=str(build_dir), TEMP=str(build_dir)) + compile_started = time.perf_counter() + returncode = None + failure = None + try: + with (build_dir / "compiler.stdout.log").open("xb") as stdout, \ + (build_dir / "compiler.stderr.log").open("xb") as stderr: + process = subprocess.run(command, cwd=build_dir, env=environment, + stdout=stdout, stderr=stderr, timeout=120, check=False) + returncode = process.returncode + except (OSError, subprocess.TimeoutExpired) as exc: + failure = f"{type(exc).__name__}: {exc}" + compile_seconds = time.perf_counter() - compile_started + complete = returncode == 0 and cubin.is_file() and cubin.stat().st_size > 0 + receipt = dict(complete=complete, identity=identity, command=command, + returncode=returncode, failure=failure, + nvcc_compile_seconds=compile_seconds, + files_sha256={path.name: _file_sha256(path) for path in sorted(build_dir.iterdir()) + if path.is_file() and not path.is_symlink()}) + _save_json_new(receipt_path, receipt) + if not complete: + raise RuntimeError("explicit nvcc candidate compilation failed; evidence retained at " + str(build_dir)) + return cubin, dict(cache_key=key, cache_directory=str(build_dir), cache_hit=False, + receipt_sha256=_file_sha256(receipt_path), **receipt), dict( + compiler_probe_seconds=probe_seconds, nvcc_compile_seconds=compile_seconds) + + def _require(condition: Any, message: str) -> None: if not condition: raise ValueError(message) @@ -316,21 +420,24 @@ def cpu_reference(prepared: PreparedEvents, candidate_sources: np.ndarray, class CudaEligibility: - """CuPy RawKernel candidate with explicit round-to-nearest float32 math. + """Explicit nvcc cubin, loaded by CuPy, with separate float32 rounding. Use a dedicated environment and set an absolute CUPY_CACHE_DIR outside any - frozen tree before construction. The first run explicitly compiles/loads both - kernels. cold_compile_seconds measures that first-use compile/cache-load, - and does not prove an empty disk cache. Construction time is reported apart. + frozen tree before construction. The first run builds or authenticates one + nvcc cubin and loads its kernels with RawModule(path=...). CuPy's compiler + is not called, so it cannot inject an incompatible FTZ option. Cold timing + separates nvcc compilation, compiler probing and driver module loading. run includes validation, transfers and output ownership in end-to-end wall. """ - def __init__(self, device_index: int = 0): + def __init__(self, device_index: int = 0, *, nvcc_path: str = DEFAULT_NVCC): started = time.perf_counter() self.device_index = _integer(device_index, "device_index", 0) cache_dir = os.environ.get("CUPY_CACHE_DIR", "") _require(bool(cache_dir) and Path(cache_dir).is_absolute(), "set CUPY_CACHE_DIR to an absolute candidate cache directory outside frozen trees") + self.cache_root = Path(cache_dir) + self.nvcc_path = Path(nvcc_path) try: import cupy as cp except ImportError as exc: @@ -339,8 +446,6 @@ class CudaEligibility: self.device = cp.cuda.Device(self.device_index) with self.device: self.stream = cp.cuda.Stream(non_blocking=True) - self.trace_kernel = cp.RawKernel(CUDA_SOURCE, "eligibility_trace", options=CUDA_OPTIONS) - self.delta_kernel = cp.RawKernel(CUDA_SOURCE, "eligibility_delta", options=CUDA_OPTIONS) self.device.synchronize() self.initialization_seconds = time.perf_counter() - started self._compiled = False @@ -366,12 +471,20 @@ class CudaEligibility: cp = self.cp first_run = not self._compiled compile_seconds = 0.0 + compilation_timing = dict(compiler_probe_seconds=0.0, nvcc_compile_seconds=0.0, + module_load_seconds=0.0) with self.device, self.stream: if first_run: compile_started = time.perf_counter() - self.trace_kernel.compile() - self.delta_kernel.compile() + cubin, self.compilation_audit, probe_timing = _nvcc_cubin( + self.cache_root, str(self.device.compute_capability), self.nvcc_path) + compilation_timing.update(probe_timing) + load_started = time.perf_counter() + self.module = cp.RawModule(path=str(cubin)) + self.trace_kernel = self.module.get_function("eligibility_trace") + self.delta_kernel = self.module.get_function("eligibility_delta") self.stream.synchronize() + compilation_timing["module_load_seconds"] = time.perf_counter() - load_started compile_seconds = time.perf_counter() - compile_started self._compiled = True transfer_started = time.perf_counter() @@ -438,7 +551,8 @@ class CudaEligibility: guard_seconds = time.perf_counter() - guard_started timing = dict(first_run=first_run, initialization_seconds=self.initialization_seconds, cold_compile_seconds=compile_seconds, - compile_timing_semantics="first-use compile/cache-load; cache coldness externally controlled", + **compilation_timing, + compile_timing_semantics="first-use nvcc probe/build or authenticated cache, then driver module load", preparation_seconds=prepared.preparation_seconds, run_validation_seconds=validation_seconds, host_to_device_and_allocation_seconds=upload_seconds, @@ -450,9 +564,11 @@ class CudaEligibility: kernel_statistic="median of resident timed repeats", end_to_end_semantics="all requested warmups/repeats and transfers; preparation excluded", end_to_end_seconds=time.perf_counter() - started) - audit = _audit(prepared, work, "cupy-rawkernel-independent-candidate") + audit = _audit(prepared, work, "nvcc-cubin-cupy-independent-candidate") audit.update(numpy_version=str(np.__version__), cupy_version=str(cp.__version__), device_index=self.device_index, cuda_options=list(CUDA_OPTIONS), + compiler_backend="external nvcc --cubin; CuPy RawModule path loading only", + compilation=self.compilation_audit, float32_separate_rounding_intrinsics=True, tf32_used=False, numerical_equivalence_established_by_this_call=False) return EligibilityResult(sampled, final, snapshots, output, timing, audit) diff --git a/models/male-cns-lif/conditioning-acceleration-v1/test_gpu_math.py b/models/male-cns-lif/conditioning-acceleration-v1/test_gpu_math.py index 24a44ce..08d52c1 100644 --- a/models/male-cns-lif/conditioning-acceleration-v1/test_gpu_math.py +++ b/models/male-cns-lif/conditioning-acceleration-v1/test_gpu_math.py @@ -6,12 +6,15 @@ import importlib.util import math import os from pathlib import Path +import subprocess import sys +import tempfile import unittest +from unittest import mock import numpy as np -from gpu_backend import (CUDA_OPTIONS, CUDA_SOURCE, COMPARTMENTS, CudaEligibility, cpu_reference, +from gpu_backend import (CUDA_OPTIONS, CUDA_SOURCE, COMPARTMENTS, CudaEligibility, _nvcc_cubin, cpu_reference, prepare_events) @@ -229,5 +232,84 @@ class HostMathTests(unittest.TestCase): self.assertNotIn("__expf", CUDA_SOURCE) +class CompilerCacheTests(unittest.TestCase): + """Host-side compiler invocation/evidence checks; these do not compile CUDA.""" + + def test_explicit_flags_authenticated_cache_and_no_global_environment_mutation(self): + with tempfile.TemporaryDirectory() as directory: + root = Path(directory) + compiler = root / "nvcc" + compiler.write_bytes(b"host-test-compiler-placeholder") + calls = [] + + def run(argv, **kwargs): + if argv[-1] == "--version": + return subprocess.CompletedProcess(argv, 0, stdout="nvcc test version", stderr="") + calls.append((argv, kwargs)) + self.assertIn("--ftz=false", argv) + self.assertIn("--fmad=false", argv) + self.assertIn("-arch=sm_70", argv) + self.assertNotIn("--use_fast_math", argv) + self.assertNotIn("-ftz=true", argv) + Path(argv[argv.index("-o") + 1]).write_bytes(b"test cubin, not executable CUDA") + kwargs["stdout"].write(b"mock nvcc stdout\n") + kwargs["stderr"].write(b"mock nvcc stderr\n") + return subprocess.CompletedProcess(argv, 0) + + with mock.patch.dict(os.environ, {"NVCC_PREPEND_FLAGS": "", "NVCC_APPEND_FLAGS": ""}), \ + mock.patch("gpu_backend.subprocess.run", side_effect=run): + before = dict(os.environ) + cubin, evidence, timing = _nvcc_cubin(root / "cache", "70", compiler) + self.assertFalse(evidence["cache_hit"]) + self.assertTrue(evidence["complete"]) + self.assertIn("eligibility.cubin", evidence["files_sha256"]) + self.assertIn("compiler.stdout.log", evidence["files_sha256"]) + self.assertIn("compiler.stderr.log", evidence["files_sha256"]) + self.assertEqual(before, dict(os.environ)) + again, cached, second_timing = _nvcc_cubin(root / "cache", "70", compiler) + self.assertEqual(cubin, again) + self.assertTrue(cached["cache_hit"]) + self.assertEqual(second_timing["nvcc_compile_seconds"], 0) + self.assertEqual(len(calls), 1) + self.assertEqual(evidence["receipt_sha256"], cached["receipt_sha256"]) + cubin.write_bytes(b"corrupted test cache") + with self.assertRaisesRegex(ValueError, "artifact changed"): + _nvcc_cubin(root / "cache", "70", compiler) + + def test_failed_compile_evidence_is_retained_and_not_overwritten(self): + with tempfile.TemporaryDirectory() as directory: + root = Path(directory) + compiler = root / "nvcc" + compiler.write_bytes(b"host-test-compiler-placeholder") + + def run(argv, **kwargs): + if argv[-1] == "--version": + return subprocess.CompletedProcess(argv, 0, stdout="nvcc test version", stderr="") + kwargs["stderr"].write(b"mock compiler failure\n") + return subprocess.CompletedProcess(argv, 1) + + with mock.patch.dict(os.environ, {"NVCC_PREPEND_FLAGS": "", "NVCC_APPEND_FLAGS": ""}), \ + mock.patch("gpu_backend.subprocess.run", side_effect=run): + with self.assertRaisesRegex(RuntimeError, "evidence retained"): + _nvcc_cubin(root / "cache", "70", compiler) + receipts = list((root / "cache").rglob("build-receipt.json")) + self.assertEqual(len(receipts), 1) + before = receipts[0].read_bytes() + with self.assertRaisesRegex(ValueError, "failed or mismatched"): + _nvcc_cubin(root / "cache", "70", compiler) + self.assertEqual(receipts[0].read_bytes(), before) + + def test_injected_compiler_flags_rejected_without_clearing_global_environment(self): + with tempfile.TemporaryDirectory() as directory: + root = Path(directory) + compiler = root / "nvcc" + compiler.write_bytes(b"host-test-compiler-placeholder") + with mock.patch.dict(os.environ, {"NVCC_APPEND_FLAGS": "--ftz=true"}): + with self.assertRaisesRegex(ValueError, "injected flags"): + _nvcc_cubin(root / "cache", "70", compiler) + self.assertEqual(os.environ["NVCC_APPEND_FLAGS"], "--ftz=true") + self.assertFalse((root / "cache").exists()) + + if __name__ == "__main__": unittest.main()