range / cuda /src /range_curate_cuda /_runtime /phase38_kernel.cu
Brianferrell787's picture
Publish retained Phase60 CUDA preview with matching installation wheels
8947dc2 verified
Raw History Blame Contribute Delete
24.6 kB
// Phase38 research: identical frozen preparation and row evaluator for AC/PFAC.
// Local Phase32B helpers/evaluator retained; no external library source is used.
#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;
}
struct MatchEvent { int key, start, end, stream; };
extern "C" __global__ void range38_prepare(
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, unsigned int* raw_source_offsets,
int* normalized_lengths, int* source_lengths, unsigned int* event_counts,
unsigned char* codes, unsigned int* facts, unsigned long long* stats, int collect_stats
) {
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;
normalized_lengths[row] = source_lengths[row] = 0;
event_counts[row] = 0;
if (collect_stats) for (int i = 0; i < 25; ++i) stats[(long long)row * 25 + 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;
}
}
}
for (int unit = 0; unit < width; ++unit)
raw_source_offsets[base + i + unit] = source_length;
i += width;
++source_length;
}
normalized_lengths[row] = norm_length;
source_lengths[row] = source_length;
if (collect_stats) stats[(long long)row * 25 + 23] = norm_length;
}
// Saturating reservation: cap+1 is a persistent overflow marker, never a wrap.
// No consumer reads any event until the complete matching kernel has finished.
__device__ bool range38_emit(
MatchEvent* events, unsigned int* counter, int cap, int key,
int start, int end, int stream
) {
unsigned int observed = atomicAdd(counter, 0u);
for (;;) {
if (observed > (unsigned int)cap) return false;
unsigned int desired = observed + 1u;
unsigned int prior = atomicCAS(counter, observed, desired);
if (prior == observed) {
if (observed == (unsigned int)cap) return false;
events[observed] = MatchEvent{key, start, end, stream};
return true;
}
observed = prior;
}
}
__device__ __forceinline__ bool range38_boundary(
const unsigned char* text, int n, int start, int end, int stream
) {
return (start == 0 || !(stream == 0 ? token(text[start - 1]) : exact_token(text[start - 1]))) &&
(end == n || !(stream == 0 ? token(text[end]) : exact_token(text[end])));
}
__device__ void range38_depth(unsigned long long* stats, int depth) {
atomicMax(stats + 24, (unsigned long long)depth);
int bucket = 0;
while (depth && bucket < 17) { ++bucket; depth >>= 1; }
atomicAdd(stats + 4 + bucket, 1ull);
}
// Ordinary completed AC scans each entire stream in one thread per original row.
extern "C" __global__ void range38_match_ac(
const unsigned char* data, const long long* row_offsets, int rows,
const unsigned char* normalized_data, const unsigned int* source_offsets,
const unsigned int* raw_source_offsets, const int* normalized_lengths,
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 unsigned char* codes, MatchEvent* all_events, unsigned int* event_counts,
int event_capacity, unsigned long long* all_stats, int collect_stats
) {
int row = blockIdx.x * blockDim.x + threadIdx.x;
if (row >= rows || codes[row] >= 2) return;
long long base = row_offsets[row];
MatchEvent* events = all_events + (long long)row * event_capacity;
unsigned long long* stats = all_stats + (long long)row * 25;
unsigned long long jobs = 0, attempts = 0, terminals = 0, accepted = 0;
bool full = false;
for (int stream = 0; stream < 3 && !full; ++stream) {
const int* meta = dfa_meta + stream * 4;
if (!meta[3]) continue;
const unsigned char* text = stream == 0 ? normalized_data + base * 2 : data + base;
const unsigned int* offsets = stream == 0 ? source_offsets + base * 2 : raw_source_offsets + base;
int n = stream == 0 ? normalized_lengths[row] : (int)(row_offsets[row + 1] - base);
unsigned int state = 0;
int depth = 0;
++jobs;
for (int pos = 0; pos < n && !full; ++pos) {
unsigned char c = stream == 2 ? lower(text[pos]) : text[pos];
state = transitions[meta[0] + state * meta[1] + classes[stream * 256 + c]];
++attempts; ++depth;
for (int out = output_offsets[meta[2] + state]; out < output_offsets[meta[2] + state + 1]; ++out) {
++terminals;
int key = output_keys[out], end = pos + 1, start = end - key_lengths[key];
if (!range38_boundary(text, n, start, end, stream)) continue;
++accepted;
if (!range38_emit(events, event_counts + row, event_capacity, key,
offsets[start], offsets[end - 1] + 1, stream)) { full = true; break; }
}
}
if (collect_stats) range38_depth(stats, depth);
}
if (collect_stats) {
stats[0] = jobs; stats[1] = attempts; stats[2] = terminals; stats[3] = accepted;
stats[22] = full ? 1 : 0;
}
}
// One block owns an article. Each lane independently walks a genuine trie from
// its starting byte positions; every terminal on that path is reported. Missing
// trie edges stop that start, unlike a failure-completed AC transition.
extern "C" __global__ void range38_match_pfac(
const unsigned char* data, const long long* row_offsets, int rows,
const unsigned char* normalized_data, const unsigned int* source_offsets,
const unsigned int* raw_source_offsets, const int* normalized_lengths,
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 unsigned char* codes, MatchEvent* all_events, unsigned int* event_counts,
int event_capacity, unsigned long long* all_stats, int collect_stats
) {
int row = blockIdx.x;
if (row >= rows || codes[row] >= 2) return;
long long base = row_offsets[row];
MatchEvent* events = all_events + (long long)row * event_capacity;
unsigned long long* stats = all_stats + (long long)row * 25;
unsigned long long jobs = 0, attempts = 0, terminals = 0, accepted = 0, stopped = 0;
for (int stream = 0; stream < 3; ++stream) {
const int* meta = dfa_meta + stream * 4;
if (!meta[3]) continue;
const unsigned char* text = stream == 0 ? normalized_data + base * 2 : data + base;
const unsigned int* offsets = stream == 0 ? source_offsets + base * 2 : raw_source_offsets + base;
int n = stream == 0 ? normalized_lengths[row] : (int)(row_offsets[row + 1] - base);
for (int start = threadIdx.x; start < n; start += blockDim.x) {
// Reserve atomically only when a match is emitted. Polling the one
// row counter for every starting byte serializes unrelated misses.
// The saturating emitter stops this lane on its next overflowing hit.
unsigned int state = 0;
int depth = 0;
bool full = false;
++jobs;
for (int pos = start; pos < n; ++pos) {
unsigned char c = stream == 2 ? lower(text[pos]) : text[pos];
state = transitions[meta[0] + state * meta[1] + classes[stream * 256 + c]];
++attempts; ++depth;
if (!state) break;
for (int out = output_offsets[meta[2] + state]; out < output_offsets[meta[2] + state + 1]; ++out) {
++terminals;
int key = output_keys[out], end = pos + 1;
if (!range38_boundary(text, n, start, end, stream)) continue;
++accepted;
if (!range38_emit(events, event_counts + row, event_capacity, key,
offsets[start], offsets[end - 1] + 1, stream)) { full = true; break; }
}
if (full) break;
}
if (collect_stats) range38_depth(stats, depth);
if (full) { ++stopped; break; }
}
}
if (collect_stats) {
atomicAdd(stats, jobs); atomicAdd(stats + 1, attempts);
atomicAdd(stats + 2, terminals); atomicAdd(stats + 3, accepted);
atomicAdd(stats + 22, stopped);
}
}
extern "C" __global__ void range38_evaluate(
const long long* row_offsets, const long long* timestamps, int rows,
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, int conjunction_count,
const unsigned char* normalized_data, const unsigned int* source_offsets,
const int* normalized_lengths, const int* source_lengths,
const MatchEvent* all_events, const unsigned int* event_counts, int event_capacity,
Candidate* all_candidates, unsigned char* codes, unsigned int* facts
) {
int row = blockIdx.x * blockDim.x + threadIdx.x;
if (row >= rows || codes[row] >= 2) return;
unsigned int* fact = facts + (long long)row * 8;
if (event_counts[row] > (unsigned int)event_capacity) { codes[row] = 3; return; }
long long base = row_offsets[row];
const unsigned char* normalized = normalized_data + base * 2;
const 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;
const MatchEvent* events = all_events + (long long)row * event_capacity;
int count = 0;
unsigned char conjunctions[MAX_CONJUNCTIONS] = {0};
bool native_candidate = false, exact_candidate = false;
for (unsigned int event = 0; event < event_counts[row]; ++event) {
MatchEvent value = events[event];
if (value.stream == 0) native_candidate = true; else exact_candidate = true;
if (!append_hit(value.key, value.start, value.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;
}