Spaces:
Running on Zero
Running on Zero
| // Phase40 research: this separately bound include supplies frozen preparation, | |
| // original AC/PFAC controls, event reservation and the exact row evaluator. | |
| extern "C" __global__ void range40_match_segmented( | |
| 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 int* stream_max_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 tiles = 0, attempts = 0, terminals = 0, accepted = 0, halo = 0; | |
| bool lane_full = false; | |
| for (int stream = 0; stream < 3 && !lane_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); | |
| int tile_count = (n + PHASE40_TILE_BYTES - 1) / PHASE40_TILE_BYTES; | |
| int maximum_length = stream_max_lengths[stream]; | |
| for (int tile = threadIdx.x; tile < tile_count && !lane_full; tile += blockDim.x) { | |
| int owned_start = tile * PHASE40_TILE_BYTES; | |
| int owned_end = owned_start + PHASE40_TILE_BYTES; | |
| if (owned_end > n) owned_end = n; | |
| int warm_start = owned_start - (maximum_length - 1); | |
| if (warm_start < 0) warm_start = 0; | |
| unsigned int state = 0; | |
| int depth = 0; | |
| ++tiles; | |
| // An ending byte in the owned tile cannot need more than Lmax-1 | |
| // preceding bytes. Starting AC at root here therefore recovers all | |
| // owned hits, even if this byte lies inside a UTF-8 encoding. | |
| for (int pos = warm_start; pos < owned_end && !lane_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; | |
| if (pos < owned_start) { ++halo; continue; } | |
| // Suppress warm-up outputs. Exactly one tile owns each final | |
| // byte, so overlapping tiles never duplicate a key occurrence. | |
| 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]; | |
| // Boundary tests use whole-row neighbors, not tile edges. | |
| 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)) { lane_full = true; break; } | |
| } | |
| } | |
| if (collect_stats) range38_depth(stats, depth); | |
| } | |
| } | |
| if (collect_stats) { | |
| atomicAdd(stats, tiles); atomicAdd(stats + 1, attempts); | |
| atomicAdd(stats + 2, terminals); atomicAdd(stats + 3, accepted); | |
| // Phase40 diagnostic slot22 is halo work; event_counts/codes separately | |
| // expose bounded-output overflow, including partially completed tiles. | |
| atomicAdd(stats + 22, halo); | |
| } | |
| } | |