Spaces:
Running on Zero
Running on Zero
| 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; | |
| } | |
| extern "C" __global__ void range_frozen_unicode_classify( | |
| 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 | |
| ) { | |
| int row = blockIdx.x * blockDim.x + threadIdx.x; | |
| if (row >= rows) return; | |
| unsigned int* fact = facts + (long long)row * 8; | |
| for (int i = 0; i < 8; ++i) fact[i] = 0; | |
| codes[row] = 0; // 2=original raw4, 3=candidate overflow, 4=oversize, 5=invalid input UTF-8, 6=2x capacity, 7=frozen table guard. | |
| long long base = row_offsets[row], raw_length = row_offsets[row + 1] - base; | |
| if (raw_length > max_row_bytes) { codes[row] = 4; return; } | |
| int length = (int)raw_length; | |
| const unsigned char* raw = data + base; | |
| // Raw row offsets partition the doubled arena: no row borrows another's space. | |
| unsigned char* normalized = normalized_data + base * 2; | |
| unsigned int* offsets = source_offsets + base * 2; | |
| int norm_capacity = length * 2; | |
| 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) { codes[row] = 5; return; } | |
| // 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) { | |
| codes[row] = 2; return; | |
| } | |
| if (scalar < 128) { | |
| if (!phase32b_emit_scalar(normalize((unsigned char)scalar), source_length, | |
| python_whitespace, normalized, offsets, norm_capacity, &norm_length, &pending_space)) { | |
| codes[row] = 6; return; | |
| } | |
| } 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)) { | |
| codes[row] = 6; return; | |
| } | |
| } 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) { codes[row] = 7; return; } | |
| int mapped_bytes = (int)unicode_pool[mapping] | ((int)unicode_pool[mapping + 1] << 8); | |
| if ((unsigned int)mapped_bytes > PHASE32B_POOL_BYTES - mapping - 2u) { | |
| codes[row] = 7; return; | |
| } | |
| 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) { codes[row] = 7; return; } | |
| if (!phase32b_emit_scalar(mapped_scalar, source_length, python_whitespace, | |
| normalized, offsets, norm_capacity, &norm_length, &pending_space)) { | |
| codes[row] = 6; return; | |
| } | |
| position += mapped_width; | |
| } | |
| } | |
| } | |
| i += width; | |
| ++source_length; | |
| } | |
| 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; | |
| } | |