// 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);
    }
}