Spaces:
Running on Zero
Running on Zero
| """Phase32B full frozen Unicode content with exact raw4/capacity CPU fallback. | |
| Taxonomy support remains the same ASCII-literal subset as A. | |
| The standalone CUDA file is read for build/source identity; runtime loads an AOT cubin. | |
| """ | |
| from __future__ import annotations | |
| import hashlib | |
| from pathlib import Path | |
| from typing import Callable | |
| import numpy as np | |
| from . import phase32a_gpu | |
| from .phase32a_table_builder import read_frozen, SOURCE_SHA256 | |
| _SOURCE_SHA256_AT_IMPORT = hashlib.sha256(Path(__file__).read_bytes()).hexdigest() | |
| _BASE_SOURCE_SHA256_AT_IMPORT = hashlib.sha256(Path(phase32a_gpu.__file__).read_bytes()).hexdigest() | |
| KERNEL_SOURCE_PATH = Path(__file__).with_name("phase32b_kernel.cu") | |
| CUDA_SOURCE = KERNEL_SOURCE_PATH.read_text(encoding="utf-8") | |
| _KERNEL_FILE_SHA256_AT_IMPORT = hashlib.sha256(KERNEL_SOURCE_PATH.read_bytes()).hexdigest() | |
| MAX_CANDIDATES = phase32a_gpu.MAX_CANDIDATES | |
| NORMALIZATION_CAPACITY_MULTIPLIER = 2 | |
| PHASE32B_SOURCE_SHA256 = SOURCE_SHA256 | |
| def _load_frozen_tables(directory): | |
| assets, meta = read_frozen(directory) | |
| arrays = { | |
| "unicode_index": np.frombuffer(assets["unicode_index.bin"], dtype="<u4"), | |
| "unicode_pool": np.frombuffer(assets["unicode_pool.bin"], dtype=np.uint8), | |
| "python_whitespace": np.frombuffer(assets["python_whitespace.bin"], dtype=np.uint8), | |
| } | |
| return arrays, { | |
| "predicate_id": "range-u13-full-content-exclude-raw4-2x-v1", | |
| "source_sha256": SOURCE_SHA256, | |
| "source_metadata": meta, | |
| "payloads": {name: {"bytes": len(data), "sha256": SOURCE_SHA256[name]} | |
| for name, data in assets.items()}, | |
| "mapping_storage": "unchanged flat frozen index, pool and whitespace assets", | |
| "normalization_capacity_multiplier": 2, | |
| "raw_excluded_scalars": [0x0130, 0x0131, 0x017f, 0x212a], | |
| } | |
| class ExactFrozenUnicodeCudaClassifier(phase32a_gpu.ExactSelectiveUnicodeCudaClassifier): | |
| def __init__(self, cpu_engine, cubin_path, *, masked_cpu: Callable, unicode_tables, | |
| batch_bytes=64 * 1024 * 1024, batch_rows=65536, max_row_bytes=65536, | |
| overlap_cpu=False, diagnostics=False): | |
| from range_curate._core.accelerated.compact import _build_native_specs | |
| for name, value in (("batch_bytes", batch_bytes), ("batch_rows", batch_rows), | |
| ("max_row_bytes", max_row_bytes)): | |
| if isinstance(value, bool) or not isinstance(value, int) or not 0 < value < 2**31: | |
| raise ValueError(f"{name} must be a positive int32") | |
| if batch_bytes > 64 * 1024 * 1024 or batch_rows > 65536 or max_row_bytes > 65536: | |
| raise ValueError("Phase32B initial caps are 64 MiB / 65,536 rows / 65,536 row bytes") | |
| if max_row_bytes > batch_bytes: | |
| raise ValueError("max_row_bytes must not exceed batch_bytes") | |
| if not callable(masked_cpu): | |
| raise TypeError("masked_cpu must be the current native masked scanner callable") | |
| if not isinstance(overlap_cpu, bool): | |
| raise TypeError("overlap_cpu must be a boolean") | |
| if not isinstance(diagnostics, bool): | |
| raise TypeError("diagnostics must be a boolean") | |
| self.cpu_engine, self.masked_cpu = cpu_engine, masked_cpu | |
| self.cubin_path = Path(cubin_path).resolve() | |
| self.batch_bytes, self.batch_rows, self.max_row_bytes = batch_bytes, batch_rows, max_row_bytes | |
| self.overlap_cpu = overlap_cpu | |
| self.diagnostics = diagnostics | |
| self._unicode_arrays, self._unicode_identity = _load_frozen_tables(unicode_tables) | |
| self.unsupported_reason = None | |
| self._arrays, self._conjunction_count = {}, 0 | |
| try: | |
| self._arrays, self._conjunction_count = phase32a_gpu._pack_specs(_build_native_specs(cpu_engine._screening_indicator)) | |
| except phase32a_gpu.UnsupportedGpuTaxonomy as exc: | |
| self.unsupported_reason = str(exc) | |
| if not self.unsupported_reason: | |
| self._arrays.update(self._unicode_arrays) | |
| self.cp = self._kernel = None | |
| self._device_config, self._workspace = {}, {} | |
| self._capacity_bytes = self._capacity_rows = 0 | |
| self._cubin_sha256 = None | |
| self._cuda_source_sha256 = hashlib.sha256(CUDA_SOURCE.encode()).hexdigest() | |
| self._kernel_attributes = None | |
| self.last_profile = {} | |
| def backend_report(self): | |
| return {**super().backend_report, | |
| "backend": "cpu" if self.unsupported_reason else "experimental_frozen_unicode_cuda", | |
| "unicode_fallback_definition": "original raw4; code6 separately records whole-row 2x capacity rejection", | |
| "taxonomy_scope": "unchanged A ASCII pattern, context and exact-literal support", | |
| "normalization_capacity_multiplier": 2, | |
| "phase32b_kernel_file_sha256": _KERNEL_FILE_SHA256_AT_IMPORT} | |
| def _initialize(self): | |
| if self.cp is not None: | |
| return | |
| if not self.cubin_path.is_file(): | |
| raise FileNotFoundError("Build the experimental cubin explicitly before running CUDA: " + str(self.cubin_path)) | |
| import cupy as cp | |
| module = cp.RawModule(path=str(self.cubin_path)) | |
| kernel = module.get_function("range_frozen_unicode_classify") | |
| # Keep load-time metadata: later reports must not describe a replacement | |
| # file at the same path or rehash source text inside screening timings. | |
| loaded_cubin_sha256 = hashlib.sha256(self.cubin_path.read_bytes()).hexdigest() | |
| attributes = kernel.attributes | |
| kernel_attributes = { | |
| name: int(attributes[name]) if attributes.get(name) is not None else None | |
| for name in ("num_regs", "local_size_bytes", "shared_size_bytes") | |
| } | |
| config = {name: cp.asarray(value if value.size else np.zeros(1, dtype=value.dtype)) | |
| for name, value in self._arrays.items()} | |
| cp.cuda.Stream.null.synchronize() | |
| self.cp, self._module, self._kernel, self._device_config = cp, module, kernel, config | |
| self._cubin_sha256 = loaded_cubin_sha256 | |
| self._kernel_attributes = kernel_attributes | |
| def _reserve(self, byte_count, row_count): | |
| cp = self.cp | |
| if byte_count > self._capacity_bytes: | |
| self._capacity_bytes = max(byte_count, self._capacity_bytes) | |
| self._workspace.update(data=cp.empty(self._capacity_bytes, dtype=cp.uint8), | |
| normalized=cp.empty(2 * self._capacity_bytes, dtype=cp.uint8), | |
| source_offsets=cp.empty(2 * self._capacity_bytes, dtype=cp.uint32)) | |
| if row_count > self._capacity_rows: | |
| self._capacity_rows = max(row_count, self._capacity_rows) | |
| self._workspace.update(offsets=cp.empty(self._capacity_rows + 1, dtype=cp.int64), | |
| timestamps=cp.empty(self._capacity_rows, dtype=cp.int64), | |
| candidates=cp.empty(self._capacity_rows * MAX_CANDIDATES * 4, dtype=cp.int32), | |
| codes=cp.empty(self._capacity_rows, dtype=cp.uint8), | |
| facts=cp.empty((self._capacity_rows, 8), dtype=cp.uint32)) | |
| def _gpu_stage(self, batch, *, text_column, timestamp_column, diagnostic=None): | |
| stage = super()._gpu_stage(batch, text_column=text_column, | |
| timestamp_column=timestamp_column, diagnostic=diagnostic) | |
| # Frozen payload hashes make code7 impossible in valid operation. Refuse | |
| # to continue or export a corrupted-table result to the native scanner. | |
| if np.any(stage[0] == 7): | |
| raise RuntimeError("Phase32B frozen mapping guard failed") | |
| return stage | |
| def diagnostic_unicode_shadow(self, data, *, enabled=False, text_column="text"): | |
| raise NotImplementedError("Phase32B needs a separate bounded full-repertoire shadow; do not use A's subset model") | |