// Phase52: warp-parallel frozen normalization, then original B matching/evaluation. // All exceptional rows replay the unchanged serial normalizer on one GPU lane. // Input/output semantics, 2x arenas, raw streams and taxonomy packing retain B. #define MAX_CANDIDATES 128 #define MAX_CONJUNCTIONS 32 #define PHASE32B_POOL_BYTES 29973u struct Candidate { int kind, index, start, end; }; __device__ __forceinline__ bool token(unsigned char c) { return (c >= 'a' && c <= 'z') || (c >= '0' && c <= '9') || c == '_'; } __device__ __forceinline__ bool exact_token(unsigned char c) { return token(c) || (c >= 'A' && c <= 'Z'); } __device__ __forceinline__ unsigned char lower(unsigned char c) { return c >= 'A' && c <= 'Z' ? c + 32 : c; } __device__ __forceinline__ unsigned char normalize(unsigned char c) { c = lower(c); if (c == '-' || c == '.' || c == ',' || c == ':' || c == ';' || c == '(' || c == ')' || c == '[' || c == ']' || c == '{' || c == '}' || c == '/') return ' '; return c == '`' ? '\'' : c; } __device__ __forceinline__ bool space(unsigned char c) { return c == ' ' || (c >= 9 && c <= 13) || (c >= 0x1c && c <= 0x1f); } // All candidate identities are unique before the fixed-size buffer is used. // An overflow transfers the whole row to CPU; no partial result is retained. __device__ bool append_hit( int key, int start, int end, const int* key_offsets, const int* key_consumers, const int* consumers, Candidate* candidates, int* count, unsigned char* conjunctions ) { for (int i = key_offsets[key]; i < key_offsets[key + 1]; ++i) { const int* spec = consumers + key_consumers[i] * 4; if (spec[0] == 2) { conjunctions[spec[1]] |= (unsigned char)(1u << spec[2]); continue; } bool duplicate = false; for (int j = 0; j < *count; ++j) { Candidate prior = candidates[j]; if (prior.kind == spec[0] && prior.index == spec[1] && prior.start == start && prior.end == end) { duplicate = true; break; } } if (duplicate) continue; if (*count == MAX_CANDIDATES) return false; candidates[*count] = Candidate{spec[0], spec[1], start, end}; ++*count; } return true; } __device__ int lower_bound_offset(const unsigned int* offsets, int n, int value) { int left = 0, right = n; while (left < right) { int middle = left + (right - left) / 2; if (offsets[middle] < (unsigned int)value) left = middle + 1; else right = middle; } return left; } __device__ bool contains_term( const unsigned char* text, int left, int right, int term, const unsigned char* term_bytes, const int* term_offsets ) { int first = term_offsets[term], length = term_offsets[term + 1] - first; if (!length || right - left < length) return false; for (int start = left; start + length <= right; ++start) { if (start > left && token(text[start - 1])) continue; if (start + length < right && token(text[start + length])) continue; int i = 0; while (i < length && text[start + i] == term_bytes[first + i]) ++i; if (i == length) return true; } return false; } __device__ void context_bounds( const unsigned char* normalized, const unsigned int* offsets, int norm_length, int source_length, int start, int end, int width, int* left, int* right ) { int source_left = start > width ? start - width : 0; int source_right = width > source_length - end ? source_length : end + width; *left = lower_bound_offset(offsets, norm_length, source_left); *right = lower_bound_offset(offsets, norm_length, source_right); // Match Phase24's borrowed-window treatment of collapsed leading spaces. if (*left < *right && normalized[*left] == ' ') ++*left; } __device__ bool alias_context_ok( int context, const int* contexts, const int* clauses, const int* term_ids, const unsigned char* term_bytes, const int* term_offsets, const unsigned char* normalized, const unsigned int* offsets, int norm_length, int source_length, int start, int end ) { if (context < 0) return true; const int* spec = contexts + context * 5; int left, right; context_bounds(normalized, offsets, norm_length, source_length, start, end, spec[0], &left, &right); for (int i = 0; i < spec[4]; ++i) if (contains_term(normalized, left, right, term_ids[spec[3] + i], term_bytes, term_offsets)) return false; for (int i = 0; i < spec[2]; ++i) { const int* clause = clauses + 2 * (spec[1] + i); bool accepted = true; for (int j = 0; j < clause[1]; ++j) { if (!contains_term(normalized, left, right, term_ids[clause[0] + j], term_bytes, term_offsets)) { accepted = false; break; } } if (accepted) return true; } return false; } // Strict UTF-8 scalar decoder. Zero width rejects the entire row. __device__ __forceinline__ int phase32b_decode( const unsigned char* raw, int length, int position, unsigned int* scalar ) { unsigned int first = raw[position]; if (first < 0x80) { *scalar = first; return 1; } int width; unsigned int value, minimum; if (first >= 0xc2 && first <= 0xdf) { width = 2; value = first & 31; minimum = 0x80; } else if (first >= 0xe0 && first <= 0xef) { width = 3; value = first & 15; minimum = 0x800; } else if (first >= 0xf0 && first <= 0xf4) { width = 4; value = first & 7; minimum = 0x10000; } else return 0; if (width > length - position) return 0; for (int index = 1; index < width; ++index) { unsigned int next = raw[position + index]; if ((next & 0xc0) != 0x80) return 0; value = (value << 6) | (next & 63); } if (value < minimum || value > 0x10ffff || (value >= 0xd800 && value <= 0xdfff)) return 0; *scalar = value; return width; } // Frozen U13 mapping output is decoded into scalars before RANGE postprocessing. // Directly copying pool bytes would skip punctuation and whitespace semantics. __device__ bool phase32b_emit_scalar( unsigned int scalar, int source_index, const unsigned char* whitespace, unsigned char* normalized, unsigned int* offsets, int capacity, int* norm_length, bool* pending_space ) { switch (scalar) { case 0x2d: case 0x2010: case 0x2011: case 0x2012: case 0x2013: case 0x2014: case 0x2015: case 0x2212: case 0x2e: case 0x2c: case 0x3a: case 0x3b: case 0x28: case 0x29: case 0x5b: case 0x5d: case 0x7b: case 0x7d: case 0xff0f: case 0x2f: scalar = 32; break; case 0x27: case 0x2019: case 0x2018: case 0x201b: case 0x60: case 0xb4: scalar = 39; break; } bool is_space = scalar < 128 ? space((unsigned char)scalar) : ((whitespace[scalar >> 3] & (1u << (scalar & 7))) != 0); if (is_space) { *pending_space = *norm_length != 0; return true; } unsigned char encoded[4]; int width; if (scalar < 0x80) { width = 1; encoded[0] = (unsigned char)scalar; } else if (scalar < 0x800) { width = 2; encoded[0] = 0xc0 | (scalar >> 6); encoded[1] = 0x80 | (scalar & 63); } else if (scalar < 0x10000) { width = 3; encoded[0] = 0xe0 | (scalar >> 12); encoded[1] = 0x80 | ((scalar >> 6) & 63); encoded[2] = 0x80 | (scalar & 63); } else { width = 4; encoded[0] = 0xf0 | (scalar >> 18); encoded[1] = 0x80 | ((scalar >> 12) & 63); encoded[2] = 0x80 | ((scalar >> 6) & 63); encoded[3] = 0x80 | (scalar & 63); } // Check the complete scalar and optional separator before either is written. int required = width + (*pending_space ? 1 : 0); if (required > capacity - *norm_length) return false; if (*pending_space) { normalized[*norm_length] = ' '; offsets[(*norm_length)++] = source_index; *pending_space = false; } for (int unit = 0; unit < width; ++unit) { normalized[*norm_length] = encoded[unit]; offsets[(*norm_length)++] = source_index; } return true; } // Exceptional rows replay B's original ordered decoder/mapping/capacity logic. // Successful speculative tiles may already have written a prefix; replay starts // at zero and overwrites it. No speculative failing tile is ever published. __device__ unsigned char phase52_serial_normalize( const unsigned char* raw, int length, const unsigned int* unicode_index, const unsigned char* unicode_pool, const unsigned char* python_whitespace, unsigned char* normalized, unsigned int* offsets, int norm_capacity, int* output_length, int* output_source_length ) { int norm_length = 0, source_length = 0; bool pending_space = false; for (int i = 0; i < length;) { unsigned int scalar; int width = phase32b_decode(raw, length, i, &scalar); if (!width) { return 5; } // Original-text Python ignorecase semantics retain these four extras. // Keep the entire original row on CPU rather than approximate raw matching. if (scalar == 0x0130 || scalar == 0x0131 || scalar == 0x017f || scalar == 0x212a) { return 2; } if (scalar < 128) { if (!phase32b_emit_scalar(normalize((unsigned char)scalar), source_length, python_whitespace, normalized, offsets, norm_capacity, &norm_length, &pending_space)) { return 6; } } else { unsigned int mapping = unicode_index[scalar]; if (mapping == 0) { if (!phase32b_emit_scalar(scalar, source_length, python_whitespace, normalized, offsets, norm_capacity, &norm_length, &pending_space)) { return 6; } } else { // The host binds exact asset hashes. Retain device range/decoder // guards as a fail-closed corruption check, not a fallback route. if (mapping > PHASE32B_POOL_BYTES - 2u) { return 7; } int mapped_bytes = (int)unicode_pool[mapping] | ((int)unicode_pool[mapping + 1] << 8); if ((unsigned int)mapped_bytes > PHASE32B_POOL_BYTES - mapping - 2u) { return 7; } const unsigned char* mapped = unicode_pool + mapping + 2; for (int position = 0; position < mapped_bytes;) { unsigned int mapped_scalar; int mapped_width = phase32b_decode(mapped, mapped_bytes, position, &mapped_scalar); if (!mapped_width) { return 7; } if (!phase32b_emit_scalar(mapped_scalar, source_length, python_whitespace, normalized, offsets, norm_capacity, &norm_length, &pending_space)) { return 6; } position += mapped_width; } } } i += width; ++source_length; } *output_length = norm_length; *output_source_length = source_length; return 0; } // A lane owns one original scalar's post-mapping piece. Its scratch payload is // trimmed at both ends, with internal spaces collapsed. All its eventual bytes // have the same original source offset, including internal collapsed separators. // The two end flags preserve whitespace that trimming removes; they are needed // when this piece is composed with preceding/following original scalars. #define PHASE52_MAX_PIECE_BYTES 33 #define PHASE52_SCRATCH_BYTES 34 #define PHASE52_WARPS_PER_BLOCK 4 #define PHASE52_FULL_MASK 0xffffffffu __device__ __forceinline__ bool phase52_piece_scalar( unsigned int scalar, const unsigned char* whitespace, unsigned char* piece, int* length, bool* pending, bool* seen_unit, bool* leading_space, bool* trailing_space ) { switch (scalar) { case 0x2d: case 0x2010: case 0x2011: case 0x2012: case 0x2013: case 0x2014: case 0x2015: case 0x2212: case 0x2e: case 0x2c: case 0x3a: case 0x3b: case 0x28: case 0x29: case 0x5b: case 0x5d: case 0x7b: case 0x7d: case 0xff0f: case 0x2f: scalar = 32; break; case 0x27: case 0x2019: case 0x2018: case 0x201b: case 0x60: case 0xb4: scalar = 39; break; } bool is_space = scalar < 128 ? space((unsigned char)scalar) : ((whitespace[scalar >> 3] & (1u << (scalar & 7))) != 0); if (!*seen_unit) { *seen_unit = true; *leading_space = is_space; } *trailing_space = is_space; if (is_space) { *pending = *length != 0; return true; } unsigned char encoded[4]; int width; if (scalar < 0x80) { width = 1; encoded[0] = (unsigned char)scalar; } else if (scalar < 0x800) { width = 2; encoded[0] = 0xc0 | (scalar >> 6); encoded[1] = 0x80 | (scalar & 63); } else if (scalar < 0x10000) { width = 3; encoded[0] = 0xe0 | (scalar >> 12); encoded[1] = 0x80 | ((scalar >> 6) & 63); encoded[2] = 0x80 | (scalar & 63); } else { width = 4; encoded[0] = 0xf0 | (scalar >> 18); encoded[1] = 0x80 | ((scalar >> 12) & 63); encoded[2] = 0x80 | ((scalar >> 6) & 63); encoded[3] = 0x80 | (scalar & 63); } if (width + (*pending ? 1 : 0) > PHASE52_MAX_PIECE_BYTES - *length) return false; if (*pending) { piece[(*length)++ * 32] = ' '; *pending = false; } for (int unit = 0; unit < width; ++unit) piece[(*length)++ * 32] = encoded[unit]; return true; } extern "C" __global__ void range52_normalize_warp( const unsigned char* data, const long long* row_offsets, int rows, int max_row_bytes, const unsigned int* unicode_index, const unsigned char* unicode_pool, const unsigned char* python_whitespace, unsigned char* normalized_data, unsigned int* source_offsets, int* normalized_lengths, int* source_lengths, unsigned char* codes, unsigned int* facts ) { // The host launch contract is exactly 128 threads: four independent warps. __shared__ unsigned char shared_payload[PHASE52_WARPS_PER_BLOCK * PHASE52_SCRATCH_BYTES * 32]; int lane = threadIdx.x & 31, local_warp = threadIdx.x >> 5; int row = blockIdx.x * PHASE52_WARPS_PER_BLOCK + local_warp; if (row >= rows) return; // Uniform for all lanes of this warp. unsigned char* piece = shared_payload + local_warp * PHASE52_SCRATCH_BYTES * 32 + lane; unsigned int* fact = facts + (long long)row * 8; if (lane < 8) fact[lane] = 0; if (lane == 0) { codes[row] = 0; normalized_lengths[row] = source_lengths[row] = 0; } long long base = row_offsets[row], raw_length = row_offsets[row + 1] - base; if (raw_length > max_row_bytes) { if (lane == 0) codes[row] = 4; return; } int length = (int)raw_length, norm_capacity = length * 2; const unsigned char* raw = data + base; unsigned char* normalized = normalized_data + base * 2; unsigned int* offsets = source_offsets + base * 2; int norm_length = 0, source_length = 0; bool pending_space = false; for (int tile = 0; tile < length; tile += 32) { int position = tile + lane, piece_length = 0; bool replay = false, valid_lead = false; bool local_pending = false, seen_unit = false, leading_space = false, trailing_space = false; if (position < length) { unsigned char first = raw[position]; if ((first & 0xc0) == 0x80) { // A continuation is an identity piece only when a strict valid // lead in the previous 1..3 bytes owns it. This also checks the // first lanes of a tile whose scalar began in the previous tile. bool covered = false; for (int back = 1; back <= 3 && back <= position; ++back) { if ((raw[position - back] & 0xc0) == 0x80) continue; unsigned int ignored; int width = phase32b_decode(raw, length, position - back, &ignored); if (width > back) { covered = true; break; } } if (!covered) replay = true; } else { unsigned int scalar; int width = phase32b_decode(raw, length, position, &scalar); if (!width) replay = true; else { valid_lead = true; if (scalar == 0x0130 || scalar == 0x0131 || scalar == 0x017f || scalar == 0x212a) replay = true; else if (scalar < 128) { replay = !phase52_piece_scalar(normalize((unsigned char)scalar), python_whitespace, piece, &piece_length, &local_pending, &seen_unit, &leading_space, &trailing_space); } else { unsigned int mapping = unicode_index[scalar]; if (!mapping) { replay = !phase52_piece_scalar(scalar, python_whitespace, piece, &piece_length, &local_pending, &seen_unit, &leading_space, &trailing_space); } else if (mapping > PHASE32B_POOL_BYTES - 2u) replay = true; else { int mapped_bytes = (int)unicode_pool[mapping] | ((int)unicode_pool[mapping + 1] << 8); if ((unsigned int)mapped_bytes > PHASE32B_POOL_BYTES - mapping - 2u || mapped_bytes > PHASE52_MAX_PIECE_BYTES) replay = true; else { const unsigned char* mapped = unicode_pool + mapping + 2; for (int at = 0; at < mapped_bytes && !replay;) { unsigned int mapped_scalar; int mapped_width = phase32b_decode(mapped, mapped_bytes, at, &mapped_scalar); if (!mapped_width) replay = true; else { replay = !phase52_piece_scalar(mapped_scalar, python_whitespace, piece, &piece_length, &local_pending, &seen_unit, &leading_space, &trailing_space); at += mapped_width; } } } } } } } } // No thread leaves before a warp-wide decision. Continuation/inactive // lanes have empty identity summaries but still participate in every scan. bool replay_row = __any_sync(PHASE52_FULL_MASK, replay); unsigned int lower_lanes = lane == 0 ? 0u : (1u << lane) - 1u; unsigned int lead_mask = __ballot_sync(PHASE52_FULL_MASK, valid_lead); unsigned int text_mask = __ballot_sync(PHASE52_FULL_MASK, piece_length != 0); unsigned int effective_mask = __ballot_sync(PHASE52_FULL_MASK, seen_unit); unsigned int effective_before = effective_mask & lower_lanes; int previous_lane = effective_before ? 31 - __clz(effective_before) : 0; bool previous_trailing = __shfl_sync(PHASE52_FULL_MASK, (int)trailing_space, previous_lane) != 0; bool incoming_space = effective_before ? previous_trailing : pending_space; bool preceding_text = norm_length != 0 || (text_mask & lower_lanes) != 0; bool separator = piece_length != 0 && preceding_text && (incoming_space || leading_space); int emitted = piece_length + (separator ? 1 : 0); int inclusive = emitted; for (int distance = 1; distance < 32; distance <<= 1) { int prior = __shfl_up_sync(PHASE52_FULL_MASK, inclusive, distance); if (lane >= distance) inclusive += prior; } int tile_bytes = __shfl_sync(PHASE52_FULL_MASK, inclusive, 31); if (tile_bytes > norm_capacity - norm_length) replay_row = true; if (replay_row) { // Serial replay resolves all ordering ambiguities, including a // capacity error earlier than a raw4 or corrupt mapping in this tile. // Speculative scratch is private and this tile has not been stored. // Order every lane's earlier global stores before lane 0 overwrites // their prefix; votes/shuffles alone are not a memory barrier. __syncwarp(PHASE52_FULL_MASK); if (lane == 0) { int replay_length = 0, replay_source_length = 0; unsigned char code = phase52_serial_normalize(raw, length, unicode_index, unicode_pool, python_whitespace, normalized, offsets, norm_capacity, &replay_length, &replay_source_length); normalized_lengths[row] = replay_length; source_lengths[row] = replay_source_length; codes[row] = code; } __syncwarp(PHASE52_FULL_MASK); return; } int original_index = source_length + __popc(lead_mask & lower_lanes); int output = norm_length + inclusive - emitted; if (separator) { normalized[output] = ' '; offsets[output++] = original_index; } for (int unit = 0; unit < piece_length; ++unit) { normalized[output] = piece[unit * 32]; offsets[output++] = original_index; } norm_length += tile_bytes; source_length += __popc(lead_mask); int final_lane = effective_mask ? 31 - __clz(effective_mask) : 0; bool final_trailing = __shfl_sync(PHASE52_FULL_MASK, (int)trailing_space, final_lane) != 0; if (effective_mask) pending_space = norm_length != 0 && final_trailing; } if (lane == 0) { normalized_lengths[row] = norm_length; source_lengths[row] = source_length; } } extern "C" __global__ void range52_match_evaluate( const unsigned char* data, const long long* row_offsets, const long long* timestamps, int rows, int max_row_bytes, const unsigned int* transitions, const unsigned short* classes, const int* dfa_meta, const int* output_offsets, const int* output_keys, const int* key_lengths, const int* key_offsets, const int* key_consumers, const int* consumers, const long long* aliases, const int* lexical, const int* contexts, const int* clauses, const int* term_ids, const unsigned char* term_bytes, const int* term_offsets, const unsigned int* unicode_index, const unsigned char* unicode_pool, const unsigned char* python_whitespace, int conjunction_count, unsigned char* normalized_data, unsigned int* source_offsets, Candidate* all_candidates, unsigned char* codes, unsigned int* facts, const int* normalized_lengths, const int* source_lengths ) { int row = blockIdx.x * blockDim.x + threadIdx.x; if (row >= rows || codes[row] >= 2) return; unsigned int* fact = facts + (long long)row * 8; long long base = row_offsets[row]; int length = (int)(row_offsets[row + 1] - base); const unsigned char* raw = data + base; unsigned char* normalized = normalized_data + base * 2; unsigned int* offsets = source_offsets + base * 2; int norm_length = normalized_lengths[row], source_length = source_lengths[row]; Candidate* candidates = all_candidates + (long long)row * MAX_CANDIDATES; int count = 0; unsigned char conjunctions[MAX_CONJUNCTIONS] = {0}; bool native_candidate = false, exact_candidate = false; // Stream 0: normalized. Stream 1: raw case-sensitive. Stream 2: raw ASCII fold. for (int stream = 0; stream < 3; ++stream) { const int* meta = dfa_meta + stream * 4; if (meta[3] == 0) continue; const unsigned char* text = stream == 0 ? normalized : raw; int n = stream == 0 ? norm_length : length; unsigned int state = 0; int raw_source_position = -1; for (int pos = 0; pos < n; ++pos) { // Validated UTF-8; every lead/ASCII byte starts one original scalar. if (stream != 0 && (text[pos] & 0xc0) != 0x80) ++raw_source_position; unsigned char c = stream == 2 ? lower(text[pos]) : text[pos]; state = transitions[meta[0] + state * meta[1] + classes[stream * 256 + c]]; int first = output_offsets[meta[2] + state]; int last = output_offsets[meta[2] + state + 1]; for (int out = first; out < last; ++out) { int key = output_keys[out]; int start = pos + 1 - key_lengths[key], end = pos + 1; bool left_ok = start == 0 || !(stream == 0 ? token(text[start - 1]) : exact_token(text[start - 1])); bool right_ok = end == n || !(stream == 0 ? token(text[end]) : exact_token(text[end])); if (!left_ok || !right_ok) continue; if (stream == 0) native_candidate = true; else exact_candidate = true; // Raw keys remain ASCII: a hit spans key_lengths[key] scalars. // Compatibility mappings never enter either raw matching stream. int original_end = stream == 0 ? offsets[end - 1] + 1 : raw_source_position + 1; int original_start = stream == 0 ? offsets[start] : original_end - key_lengths[key]; if (!append_hit(key, original_start, original_end, key_offsets, key_consumers, consumers, candidates, &count, conjunctions)) { codes[row] = 3; return; } } } } bool accepted_nonsemantic = false, accepted_evidence = false; bool unresolved_exposure = false, unresolved_event = false; unsigned int supported = 0, observations = 0; for (int i = 0; i < count; ++i) { Candidate candidate = candidates[i]; if (candidate.kind == 0) { const long long* alias = aliases + candidate.index * 8; bool longest = true; for (int j = 0; j < count; ++j) { Candidate other = candidates[j]; if (other.kind == 0 && other.start == candidate.start && aliases[other.index * 8] == alias[0] && other.end > candidate.end) { longest = false; break; } } if (!longest) continue; ++observations; // Retained observations count before temporal/context checks. if (timestamps[row] < alias[1] || (alias[3] && timestamps[row] >= alias[2])) continue; if (!alias_context_ok((int)alias[7], contexts, clauses, term_ids, term_bytes, term_offsets, normalized, offsets, norm_length, source_length, candidate.start, candidate.end)) continue; accepted_nonsemantic = true; if (alias[4] != -1 && (alias[5] == 0 || alias[5] == 1)) { accepted_evidence = true; supported |= 1u << alias[4]; } else if (alias[6]) unresolved_exposure = true; } else { const int* rule = lexical + candidate.index * 7; ++observations; bool context_ok = rule[2] == 0; if (!context_ok) { int left, right; context_bounds(normalized, offsets, norm_length, source_length, candidate.start, candidate.end, rule[0], &left, &right); for (int j = 0; j < rule[2]; ++j) { if (contains_term(normalized, left, right, term_ids[rule[1] + j], term_bytes, term_offsets)) { context_ok = true; break; } } } if (rule[3]) { if (rule[4]) { accepted_evidence = true; supported |= 8; } else { if (rule[5]) unresolved_exposure = true; if (rule[6]) unresolved_event = true; } } else { if (context_ok) accepted_nonsemantic = true; if (rule[4] && context_ok) { accepted_evidence = true; supported |= 8; } else if (rule[5]) unresolved_exposure = true; } } } for (int i = 0; i < conjunction_count; ++i) { if (conjunctions[i] == 15) { ++observations; accepted_evidence = true; supported |= 8; } } if (!observations) return; fact[0] = accepted_evidence || accepted_nonsemantic || unresolved_event ? 0 : 2; for (int family = 0; family < 4; ++family) fact[family + 1] = supported & (1u << family) ? 0 : (family != 3 && !supported && unresolved_exposure ? 1 : 2); fact[5] = observations; fact[6] = native_candidate; fact[7] = exact_candidate; codes[row] = 1; }