qishizhe/digital-drosophila
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Preserve exact float32 via explicit nvcc compilation and audited binary loading

Browse Source
This commit is contained in:
Digital Drosophila Archive
2026-10-01 14:41:59 +08:00
parent 5abb3e6463
commit 221e1f13f7
4 changed files with 326 additions and 13 deletions
Download Patch File Download Diff File Expand all files Collapse all files
models/male-cns-lif/conditioning-acceleration-v1/gpu_backend.py
+128 -12
View File
@@ -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)