Spaces:
Running on Zero
Running on Zero
File size: 4,087 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 | // Phase40 research: this separately bound include supplies frozen preparation,
// original AC/PFAC controls, event reservation and the exact row evaluator.
#include "phase38_kernel.cu"
#define PHASE40_TILE_BYTES 128
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);
}
}
|