Spaces:
Running on Zero
Running on Zero
File size: 7,996 Bytes
8947dc2 | 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59 60 61 62 63 64 65 66 67 68 69 70 71 72 73 74 75 76 77 78 79 80 81 82 83 84 85 86 87 88 89 90 91 92 93 94 95 96 97 98 99 100 101 102 103 104 105 106 107 108 109 110 111 112 113 114 115 116 117 118 119 120 121 122 123 124 125 126 127 128 129 130 131 132 133 134 135 136 137 138 139 140 141 142 143 144 145 | """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 = {}
@property
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")
|