vyre 0.1.0

/// Generate the pass-3 evaluation shader.
///
/// This shader implements the full rulefire bytecode ISA on GPU, achieving
/// CPU/interpreter parity for all 42 opcodes. Each GPU thread evaluates one
/// rule's condition program against the pre-scattered match data.
pub fn build_eval_shader(max_stack: u32, max_for_iterations: u32) -> String {
    format!(
        r#"
const STACK_SIZE: u32 = {max_stack}u;
const MAX_FOR_ITERATIONS: u32 = {max_for_iterations}u;
const MAX_LOOP_DEPTH: u32 = 8u;
const ABORT_SENTINEL: u32 = 0xFFFFFFFFu;

struct Instruction {{
    opcode: u32,
    operand: u32,
}};

struct FileContext {{
    file_size: u32,
    entropy_bucket: u32,
    magic_u32: u32,
    is_pe: u32,
    is_dll: u32,
    is_64bit: u32,
    has_signature: u32,
    num_sections: u32,
    num_imports: u32,
    entry_point_rva: u32,
    _pad0: u32,
    _pad1: u32,
}};

@group(0) @binding(0) var<storage, read> programs: array<Instruction>;
@group(0) @binding(1) var<storage, read> rule_program_spans: array<vec2<u32>>;
@group(0) @binding(2) var<storage, read> rule_bitmaps: array<u32>;
@group(0) @binding(3) var<storage, read> rule_counts: array<u32>;
@group(0) @binding(4) var<storage, read> rule_positions: array<u32>;
@group(0) @binding(5) var<uniform> params: vec4<u32>;
@group(0) @binding(6) var<storage, read> file_bytes: array<u32>;
@group(0) @binding(7) var<uniform> file_ctx: FileContext;
@group(0) @binding(8) var<storage, read> rule_lengths: array<u32>;
struct FiredResults {{
    count: atomic<u32>,
    indices: array<u32, 1024>,
}};
@group(0) @binding(9) var<storage, read_write> fired_results: FiredResults;

// params.x = rule_count
// params.y = max_strings_per_rule
// params.z = max_cached_positions
// params.w = file_size

fn count_index(rule_id: u32, string_id: u32) -> u32 {{
    return rule_id * params.y + string_id;
}}

fn pos_index(rule_id: u32, string_id: u32, slot: u32) -> u32 {{
    return ((rule_id * params.y + string_id) * params.z) + slot;
}}

// Read a single byte from the file_bytes buffer (packed as array<u32>).
fn read_byte(offset: u32) -> u32 {{
    if (offset >= file_ctx.file_size) {{
        return 0u;
    }}
    let word_idx = offset / 4u;
    if (word_idx >= arrayLength(&file_bytes)) {{
        return 0u;
    }}
    let word = file_bytes[word_idx];
    let shift = (offset % 4u) * 8u;
    return (word >> shift) & 0xFFu;
}}

// Count the number of set bits (popcount) in a u32.
fn popcount32(v: u32) -> u32 {{
    var x = v;
    x = x - ((x >> 1u) & 0x55555555u);
    x = (x & 0x33333333u) + ((x >> 2u) & 0x33333333u);
    x = (x + (x >> 4u)) & 0x0F0F0F0Fu;
    x = x + (x >> 8u);
    x = x + (x >> 16u);
    return x & 0x3Fu;
}}

// Execute a non-control-flow instruction. Returns the new stack pointer
// or ABORT_SENTINEL on stack overflow / unknown opcode.
fn exec_simple_inst(inst: Instruction, rule_id: u32, sp_in: u32, stack: ptr<function, array<u32, STACK_SIZE>>) -> u32 {{
    var sp = sp_in;
    switch inst.opcode {{
        case 1u: {{ if (sp >= STACK_SIZE) {{ return ABORT_SENTINEL; }} (*stack)[sp] = 1u; sp = sp + 1u; }}
        case 2u: {{ if (sp >= STACK_SIZE) {{ return ABORT_SENTINEL; }} (*stack)[sp] = 0u; sp = sp + 1u; }}
        case 3u: {{
            if (sp >= STACK_SIZE) {{ return ABORT_SENTINEL; }}
            let word_idx = inst.operand / 32u;
            if (word_idx >= 8u) {{ return ABORT_SENTINEL; }}
            let bit_idx = inst.operand % 32u;
            let bit = 1u << bit_idx;
            (*stack)[sp] = select(0u, 1u, (rule_bitmaps[rule_id * 8u + word_idx] & bit) != 0u);
            sp = sp + 1u;
        }}
        case 4u: {{
            if (sp >= STACK_SIZE) {{ return ABORT_SENTINEL; }}
            (*stack)[sp] = rule_counts[count_index(rule_id, inst.operand)];
            sp = sp + 1u;
        }}
        case 5u: {{
            let idx = (*stack)[sp - 1u];
            var value = 0xFFFFFFFFu; // MAX sentinel for no-match
            let count = rule_counts[count_index(rule_id, inst.operand)];
            if (idx < min(count, params.z)) {{
                value = rule_positions[pos_index(rule_id, inst.operand, idx)];
            }}
            (*stack)[sp - 1u] = value;
        }}
        case 6u: {{
            let idx = (*stack)[sp - 1u];
            var value = 0u;
            let count = rule_counts[count_index(rule_id, inst.operand)];
            if (idx < min(count, params.z)) {{
                value = rule_lengths[pos_index(rule_id, inst.operand, idx)];
            }}
            (*stack)[sp - 1u] = value;
        }}
        case 7u: {{ if (sp >= STACK_SIZE) {{ return ABORT_SENTINEL; }} (*stack)[sp] = file_ctx.file_size; sp = sp + 1u; }}
        case 8u: {{
            if (sp >= STACK_SIZE) {{ return ABORT_SENTINEL; }}
            var count = 0u;
            for (var i = 0u; i < 8u; i = i + 1u) {{
                count = count + popcount32(rule_bitmaps[rule_id * 8u + i]);
            }}
            (*stack)[sp] = count;
            sp = sp + 1u;
        }}
        case 9u: {{ if (sp >= STACK_SIZE) {{ return ABORT_SENTINEL; }} (*stack)[sp] = inst.operand; sp = sp + 1u; }}
        case 10u: {{
            if (sp >= STACK_SIZE) {{ return ABORT_SENTINEL; }}
            var count = 0u;
            for (var i = 0u; i < 8u; i = i + 1u) {{
                count = count + popcount32(rule_bitmaps[rule_id * 8u + i]);
            }}
            (*stack)[sp] = count;
            sp = sp + 1u;
        }}
        case 11u: {{ sp = sp - 1u; (*stack)[sp - 1u] = select(0u, 1u, (*stack)[sp - 1u] != 0u && (*stack)[sp] != 0u); }}
        case 12u: {{ sp = sp - 1u; (*stack)[sp - 1u] = select(0u, 1u, (*stack)[sp - 1u] != 0u || (*stack)[sp] != 0u); }}
        case 13u: {{ (*stack)[sp - 1u] = select(1u, 0u, (*stack)[sp - 1u] != 0u); }}
        case 14u: {{ sp = sp - 1u; (*stack)[sp - 1u] = select(0u, 1u, (*stack)[sp - 1u] == (*stack)[sp]); }}
        case 15u: {{ sp = sp - 1u; (*stack)[sp - 1u] = select(0u, 1u, (*stack)[sp - 1u] != (*stack)[sp]); }}
        case 16u: {{ sp = sp - 1u; (*stack)[sp - 1u] = select(0u, 1u, (*stack)[sp - 1u] < (*stack)[sp]); }}
        case 17u: {{ sp = sp - 1u; (*stack)[sp - 1u] = select(0u, 1u, (*stack)[sp - 1u] > (*stack)[sp]); }}
        case 18u: {{ sp = sp - 1u; (*stack)[sp - 1u] = select(0u, 1u, (*stack)[sp - 1u] <= (*stack)[sp]); }}
        case 19u: {{ sp = sp - 1u; (*stack)[sp - 1u] = select(0u, 1u, (*stack)[sp - 1u] >= (*stack)[sp]); }}
        case 20u: {{ sp = sp - 1u; (*stack)[sp - 1u] = (*stack)[sp - 1u] + (*stack)[sp]; }}
        case 21u: {{ sp = sp - 1u; (*stack)[sp - 1u] = (*stack)[sp - 1u] - (*stack)[sp]; }}
        case 31u: {{ sp = sp - 1u; (*stack)[sp - 1u] = (*stack)[sp - 1u] * (*stack)[sp]; }}
        case 32u: {{ sp = sp - 1u; let divisor = (*stack)[sp]; (*stack)[sp - 1u] = select((*stack)[sp - 1u] / divisor, 0u, divisor == 0u); }}
        case 33u: {{ sp = sp - 1u; let divisor = (*stack)[sp]; (*stack)[sp - 1u] = select((*stack)[sp - 1u] % divisor, 0u, divisor == 0u); }}
        case 34u: {{ sp = sp - 1u; (*stack)[sp - 1u] = (*stack)[sp - 1u] & (*stack)[sp]; }}
        case 35u: {{ sp = sp - 1u; (*stack)[sp - 1u] = (*stack)[sp - 1u] | (*stack)[sp]; }}
        case 36u: {{ sp = sp - 1u; (*stack)[sp - 1u] = (*stack)[sp - 1u] ^ (*stack)[sp]; }}
        case 37u: {{ sp = sp - 1u; let shift_amt = min((*stack)[sp], 31u); (*stack)[sp - 1u] = (*stack)[sp - 1u] << shift_amt; }}
        case 38u: {{ sp = sp - 1u; let shift_amt = min((*stack)[sp], 31u); (*stack)[sp - 1u] = (*stack)[sp - 1u] >> shift_amt; }}
        case 22u: {{
            let threshold = inst.operand & 0xFFFFu;
            let count = inst.operand >> 16u;
            var matched = 0u;
            for (var i = 0u; i < count; i = i + 1u) {{
                if ((*stack)[sp - 1u - i] != 0u) {{ matched = matched + 1u; }}
            }}
            sp = sp - count + 1u;
            (*stack)[sp - 1u] = select(0u, 1u, matched >= threshold);
        }}
        case 23u: {{
            var ok = 1u;
            for (var i = 0u; i < inst.operand; i = i + 1u) {{
                if ((*stack)[sp - 1u - i] == 0u) {{ ok = 0u; }}
            }}
            sp = sp - inst.operand + 1u;
            (*stack)[sp - 1u] = ok;
        }}
        case 24u: {{
            var ok = 0u;
            for (var i = 0u; i < inst.operand; i = i + 1u) {{
                if ((*stack)[sp - 1u - i] != 0u) {{ ok = 1u; }}
            }}
            sp = sp - inst.operand + 1u;
            (*stack)[sp - 1u] = ok;
        }}
        case 25u: {{
            let wanted = (*stack)[sp - 1u];
            let count = rule_counts[count_index(rule_id, inst.operand)];
            var ok = 0u;
            for (var i = 0u; i < min(count, params.z); i = i + 1u) {{
                if (rule_positions[pos_index(rule_id, inst.operand, i)] == wanted) {{ ok = 1u; }}
            }}
            (*stack)[sp - 1u] = ok;
        }}
        case 26u: {{
            let hi = (*stack)[sp - 1u];
            let lo = (*stack)[sp - 2u];
            let count = rule_counts[count_index(rule_id, inst.operand)];
            var ok = 0u;
            for (var i = 0u; i < min(count, params.z); i = i + 1u) {{
                let pos = rule_positions[pos_index(rule_id, inst.operand, i)];
                if (pos >= lo && pos <= hi) {{ ok = 1u; }}
            }}
            sp = sp - 1u;
            (*stack)[sp - 1u] = ok;
        }}
        case 39u: {{
            let offset = (*stack)[sp - 1u];
            let size = inst.operand & 0xFFu;
            let is_signed = ((inst.operand >> 8u) & 1u) != 0u;
            let is_be = ((inst.operand >> 9u) & 1u) != 0u;
            var value = 0u;
            if (offset + size <= file_ctx.file_size) {{
                let b0 = read_byte(offset);
                if (size == 1u) {{
                    if (is_signed && (b0 & 0x80u) != 0u) {{
                        value = b0 | 0xFFFFFF00u;
                    }} else {{
                        value = b0;
                    }}
                }} else if (size == 2u) {{
                    let b1 = read_byte(offset + 1u);
                    var v = 0u;
                    if (is_be) {{
                        v = (b0 << 8u) | b1;
                    }} else {{
                        v = (b1 << 8u) | b0;
                    }}
                    if (is_signed && (v & 0x8000u) != 0u) {{
                        value = v | 0xFFFF0000u;
                    }} else {{
                        value = v;
                    }}
                }} else if (size == 4u) {{
                    let b1 = read_byte(offset + 1u);
                    let b2 = read_byte(offset + 2u);
                    let b3 = read_byte(offset + 3u);
                    if (is_be) {{
                        value = (b0 << 24u) | (b1 << 16u) | (b2 << 8u) | b3;
                    }} else {{
                        value = (b3 << 24u) | (b2 << 16u) | (b1 << 8u) | b0;
                    }}
                }}
            }}
            (*stack)[sp - 1u] = value;
        }}
        case 41u: {{ if (sp >= STACK_SIZE) {{ return ABORT_SENTINEL; }} (*stack)[sp] = file_ctx.entropy_bucket; sp = sp + 1u; }}
        case 42u: {{ if (sp >= STACK_SIZE) {{ return ABORT_SENTINEL; }} (*stack)[sp] = file_ctx.is_pe; sp = sp + 1u; }}
        case 43u: {{ if (sp >= STACK_SIZE) {{ return ABORT_SENTINEL; }} (*stack)[sp] = file_ctx.is_dll; sp = sp + 1u; }}
        case 44u: {{ if (sp >= STACK_SIZE) {{ return ABORT_SENTINEL; }} (*stack)[sp] = file_ctx.num_sections; sp = sp + 1u; }}
        case 45u: {{ if (sp >= STACK_SIZE) {{ return ABORT_SENTINEL; }} (*stack)[sp] = file_ctx.num_imports; sp = sp + 1u; }}
        case 46u: {{ if (sp >= STACK_SIZE) {{ return ABORT_SENTINEL; }} (*stack)[sp] = file_ctx.entry_point_rva; sp = sp + 1u; }}
        case 47u: {{ if (sp >= STACK_SIZE) {{ return ABORT_SENTINEL; }} (*stack)[sp] = file_ctx.has_signature; sp = sp + 1u; }}
        case 48u: {{ if (sp >= STACK_SIZE) {{ return ABORT_SENTINEL; }} (*stack)[sp] = file_ctx.magic_u32; sp = sp + 1u; }}
        case 49u: {{ if (sp >= STACK_SIZE) {{ return ABORT_SENTINEL; }} (*stack)[sp] = file_ctx.is_64bit; sp = sp + 1u; }}
        default: {{ return ABORT_SENTINEL; }}
    }}
    return sp;
}}

// Evaluate a body (instructions from start_pc up to but not including end_pc)
// with full support for nested loops up to MAX_LOOP_DEPTH.
fn eval_body(start_pc: u32, end_pc: u32, rule_id: u32, sp_in: u32, stack: ptr<function, array<u32, STACK_SIZE>>) -> u32 {{
    if (end_pc > arrayLength(&programs)) {{
        return ABORT_SENTINEL;
    }}
    var pc = start_pc;
    var sp = sp_in;
    var depth: u32 = 0u;
    var frame_start: array<u32, 8>;
    var frame_end: array<u32, 8>;
    var frame_iter: array<u32, 8>;
    var frame_last: array<u32, 8>;
    var frame_mode: array<u32, 8>;
    var frame_n: array<u32, 8>;
    var frame_sp: array<u32, 8>;
    var frame_true_count: array<u32, 8>;

    loop {{
        if (pc >= end_pc && depth == 0u) {{ break; }}

        if (depth > 0u && pc >= frame_end[depth - 1u]) {{
            let d = depth - 1u;
            let body_result = sp > frame_sp[d] && (*stack)[sp - 1u] != 0u;
            let iter = frame_iter[d];
            let last = frame_last[d];

            if (frame_mode[d] == 0u) {{
                if (body_result) {{
                    depth = depth - 1u;
                    pc = frame_end[d] + 1u;
                    sp = frame_sp[d] + 1u;
                    (*stack)[frame_sp[d]] = 1u;
                    continue;
                }}
                if (iter >= last) {{
                    depth = depth - 1u;
                    pc = frame_end[d] + 1u;
                    sp = frame_sp[d] + 1u;
                    (*stack)[frame_sp[d]] = 0u;
                    continue;
                }}
                frame_iter[d] = iter + 1u;
                pc = frame_start[d];
                sp = frame_sp[d];
                continue;
            }}

            if (frame_mode[d] == 1u) {{
                if (!body_result) {{
                    depth = depth - 1u;
                    pc = frame_end[d] + 1u;
                    sp = frame_sp[d] + 1u;
                    (*stack)[frame_sp[d]] = 0u;
                    continue;
                }}
                if (iter >= last) {{
                    depth = depth - 1u;
                    pc = frame_end[d] + 1u;
                    sp = frame_sp[d] + 1u;
                    (*stack)[frame_sp[d]] = 1u;
                    continue;
                }}
                frame_iter[d] = iter + 1u;
                pc = frame_start[d];
                sp = frame_sp[d];
                continue;
            }}

            if (frame_mode[d] == 2u) {{
                let n = frame_n[d];
                var tc = frame_true_count[d];
                if (body_result) {{
                    tc = tc + 1u;
                    frame_true_count[d] = tc;
                }}
                if (tc >= n) {{
                    depth = depth - 1u;
                    pc = frame_end[d] + 1u;
                    sp = frame_sp[d] + 1u;
                    (*stack)[frame_sp[d]] = 1u;
                    continue;
                }}
                if (iter >= last) {{
                    depth = depth - 1u;
                    pc = frame_end[d] + 1u;
                    sp = frame_sp[d] + 1u;
                    (*stack)[frame_sp[d]] = 0u;
                    continue;
                }}
                frame_iter[d] = iter + 1u;
                pc = frame_start[d];
                sp = frame_sp[d];
                continue;
            }}
        }}

        if (pc >= end_pc) {{ break; }}

        let inst = programs[pc];
        pc = pc + 1u;
        switch inst.opcode {{
            case 27u: {{
                let body_len = inst.operand >> 16u;
                let range_upper = (*stack)[sp - 1u];
                let range_lower = (*stack)[sp - 2u];
                sp = sp - 2u;
                if (body_len > 0u && pc > 0xFFFFFFFFu - body_len) {{ return ABORT_SENTINEL; }}
                let bstart = pc;
                let bend = pc + body_len - 1u;
                pc = pc + body_len;
                if (body_len > 0u && range_lower <= range_upper) {{
                    if (depth >= MAX_LOOP_DEPTH) {{ return ABORT_SENTINEL; }}
                    frame_start[depth] = bstart;
                    frame_end[depth] = bend;
                    frame_iter[depth] = range_lower;
                    frame_last[depth] = min(range_upper, range_lower + MAX_FOR_ITERATIONS - 1u);
                    frame_mode[depth] = 0u;
                    frame_sp[depth] = sp;
                    depth = depth + 1u;
                    pc = bstart;
                }} else {{
                    if (sp >= STACK_SIZE) {{ return ABORT_SENTINEL; }}
                    (*stack)[sp] = 0u; sp = sp + 1u;
                }}
            }}
            case 28u: {{
                let body_len = inst.operand >> 16u;
                let range_upper = (*stack)[sp - 1u];
                let range_lower = (*stack)[sp - 2u];
                sp = sp - 2u;
                if (body_len > 0u && pc > 0xFFFFFFFFu - body_len) {{ return ABORT_SENTINEL; }}
                let bstart = pc;
                let bend = pc + body_len - 1u;
                pc = pc + body_len;
                if (body_len > 0u && range_lower <= range_upper) {{
                    if (depth >= MAX_LOOP_DEPTH) {{ return ABORT_SENTINEL; }}
                    frame_start[depth] = bstart;
                    frame_end[depth] = bend;
                    frame_iter[depth] = range_lower;
                    frame_last[depth] = min(range_upper, range_lower + MAX_FOR_ITERATIONS - 1u);
                    frame_mode[depth] = 1u;
                    frame_sp[depth] = sp;
                    depth = depth + 1u;
                    pc = bstart;
                }} else {{
                    if (sp >= STACK_SIZE) {{ return ABORT_SENTINEL; }}
                    (*stack)[sp] = 1u; sp = sp + 1u;
                }}
            }}
            case 40u: {{
                let n = inst.operand & 0xFFFFu;
                let body_len = inst.operand >> 16u;
                let range_upper = (*stack)[sp - 1u];
                let range_lower = (*stack)[sp - 2u];
                sp = sp - 2u;
                if (body_len > 0u && pc > 0xFFFFFFFFu - body_len) {{ return ABORT_SENTINEL; }}
                let bstart = pc;
                let bend = pc + body_len - 1u;
                pc = pc + body_len;
                if (body_len > 0u && range_lower <= range_upper) {{
                    if (depth >= MAX_LOOP_DEPTH) {{ return ABORT_SENTINEL; }}
                    frame_start[depth] = bstart;
                    frame_end[depth] = bend;
                    frame_iter[depth] = range_lower;
                    frame_last[depth] = min(range_upper, range_lower + MAX_FOR_ITERATIONS - 1u);
                    frame_mode[depth] = 2u;
                    frame_n[depth] = n;
                    frame_sp[depth] = sp;
                    frame_true_count[depth] = 0u;
                    depth = depth + 1u;
                    pc = bstart;
                }} else {{
                    if (sp >= STACK_SIZE) {{ return ABORT_SENTINEL; }}
                    (*stack)[sp] = 0u; sp = sp + 1u;
                }}
            }}
            case 29u: {{ continue; }}
            case 30u: {{ return sp; }}
            default: {{
                sp = exec_simple_inst(inst, rule_id, sp, stack);
                if (sp == ABORT_SENTINEL) {{ return ABORT_SENTINEL; }}
            }}
        }}
    }}
    return sp;
}}

@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {{
    let rule_id = gid.x;
    if (rule_id >= params.x) {{
        return;
    }}

    var stack: array<u32, {max_stack}>;
    var sp: u32 = 0u;
    let span = rule_program_spans[rule_id];
    if (span.x > 0xFFFFFFFFu - span.y) {{
        return;
    }}
    var pc = span.x;
    let end_pc = span.x + span.y;
    if (end_pc > arrayLength(&programs)) {{
        return;
    }}

    loop {{
        if (pc >= end_pc) {{
            break;
        }}
        let inst = programs[pc];
        pc = pc + 1u;
        switch inst.opcode {{
            case 1u: {{ if (sp >= STACK_SIZE) {{ return; }} stack[sp] = 1u; sp = sp + 1u; }}
            case 2u: {{ if (sp >= STACK_SIZE) {{ return; }} stack[sp] = 0u; sp = sp + 1u; }}
            case 3u: {{
                if (sp >= STACK_SIZE) {{ return; }}
                let word_idx = inst.operand / 32u;
                if (word_idx >= 8u) {{ return; }}
                let bit_idx = inst.operand % 32u;
                let bit = 1u << bit_idx;
                stack[sp] = select(0u, 1u, (rule_bitmaps[rule_id * 8u + word_idx] & bit) != 0u);
                sp = sp + 1u;
            }}
            case 4u: {{
                if (sp >= STACK_SIZE) {{ return; }}
                stack[sp] = rule_counts[count_index(rule_id, inst.operand)];
                sp = sp + 1u;
            }}
            case 5u: {{
                let idx = stack[sp - 1u];
                var value = 0xFFFFFFFFu;
                let count = rule_counts[count_index(rule_id, inst.operand)];
                if (idx < min(count, params.z)) {{
                    value = rule_positions[pos_index(rule_id, inst.operand, idx)];
                }}
                stack[sp - 1u] = value;
            }}
            case 6u: {{
                let idx = stack[sp - 1u];
                var value = 0u;
                let count = rule_counts[count_index(rule_id, inst.operand)];
                if (idx < min(count, params.z)) {{
                    value = rule_lengths[pos_index(rule_id, inst.operand, idx)];
                }}
                stack[sp - 1u] = value;
            }}
            case 7u: {{ if (sp >= STACK_SIZE) {{ return; }} stack[sp] = file_ctx.file_size; sp = sp + 1u; }}
            case 8u: {{
                if (sp >= STACK_SIZE) {{ return; }}
                var count = 0u;
                for (var i = 0u; i < 8u; i = i + 1u) {{
                    count = count + popcount32(rule_bitmaps[rule_id * 8u + i]);
                }}
                stack[sp] = count;
                sp = sp + 1u;
            }}
            case 9u: {{ if (sp >= STACK_SIZE) {{ return; }} stack[sp] = inst.operand; sp = sp + 1u; }}
            case 10u: {{
                if (sp >= STACK_SIZE) {{ return; }}
                var count = 0u;
                for (var i = 0u; i < 8u; i = i + 1u) {{
                    count = count + popcount32(rule_bitmaps[rule_id * 8u + i]);
                }}
                stack[sp] = count;
                sp = sp + 1u;
            }}
            case 11u: {{ sp = sp - 1u; stack[sp - 1u] = select(0u, 1u, stack[sp - 1u] != 0u && stack[sp] != 0u); }}
            case 12u: {{ sp = sp - 1u; stack[sp - 1u] = select(0u, 1u, stack[sp - 1u] != 0u || stack[sp] != 0u); }}
            case 13u: {{ stack[sp - 1u] = select(1u, 0u, stack[sp - 1u] != 0u); }}
            case 14u: {{ sp = sp - 1u; stack[sp - 1u] = select(0u, 1u, stack[sp - 1u] == stack[sp]); }}
            case 15u: {{ sp = sp - 1u; stack[sp - 1u] = select(0u, 1u, stack[sp - 1u] != stack[sp]); }}
            case 16u: {{ sp = sp - 1u; stack[sp - 1u] = select(0u, 1u, stack[sp - 1u] < stack[sp]); }}
            case 17u: {{ sp = sp - 1u; stack[sp - 1u] = select(0u, 1u, stack[sp - 1u] > stack[sp]); }}
            case 18u: {{ sp = sp - 1u; stack[sp - 1u] = select(0u, 1u, stack[sp - 1u] <= stack[sp]); }}
            case 19u: {{ sp = sp - 1u; stack[sp - 1u] = select(0u, 1u, stack[sp - 1u] >= stack[sp]); }}
            case 20u: {{ sp = sp - 1u; stack[sp - 1u] = stack[sp - 1u] + stack[sp]; }}
            case 21u: {{ sp = sp - 1u; stack[sp - 1u] = stack[sp - 1u] - stack[sp]; }}
            case 31u: {{ sp = sp - 1u; stack[sp - 1u] = stack[sp - 1u] * stack[sp]; }}
            case 32u: {{ sp = sp - 1u; let divisor = stack[sp]; stack[sp - 1u] = select(stack[sp - 1u] / divisor, 0u, divisor == 0u); }}
            case 33u: {{ sp = sp - 1u; let divisor = stack[sp]; stack[sp - 1u] = select(stack[sp - 1u] % divisor, 0u, divisor == 0u); }}
            case 34u: {{ sp = sp - 1u; stack[sp - 1u] = stack[sp - 1u] & stack[sp]; }}
            case 35u: {{ sp = sp - 1u; stack[sp - 1u] = stack[sp - 1u] | stack[sp]; }}
            case 36u: {{ sp = sp - 1u; stack[sp - 1u] = stack[sp - 1u] ^ stack[sp]; }}
            case 37u: {{ sp = sp - 1u; let shift_amt = min(stack[sp], 31u); stack[sp - 1u] = stack[sp - 1u] << shift_amt; }}
            case 38u: {{ sp = sp - 1u; let shift_amt = min(stack[sp], 31u); stack[sp - 1u] = stack[sp - 1u] >> shift_amt; }}
            case 22u: {{
                let threshold = inst.operand & 0xFFFFu;
                let count = inst.operand >> 16u;
                var matched = 0u;
                for (var i = 0u; i < count; i = i + 1u) {{
                    if (stack[sp - 1u - i] != 0u) {{ matched = matched + 1u; }}
                }}
                sp = sp - count + 1u;
                stack[sp - 1u] = select(0u, 1u, matched >= threshold);
            }}
            case 23u: {{
                var ok = 1u;
                for (var i = 0u; i < inst.operand; i = i + 1u) {{
                    if (stack[sp - 1u - i] == 0u) {{ ok = 0u; }}
                }}
                sp = sp - inst.operand + 1u;
                stack[sp - 1u] = ok;
            }}
            case 24u: {{
                var ok = 0u;
                for (var i = 0u; i < inst.operand; i = i + 1u) {{
                    if (stack[sp - 1u - i] != 0u) {{ ok = 1u; }}
                }}
                sp = sp - inst.operand + 1u;
                stack[sp - 1u] = ok;
            }}
            case 25u: {{
                let wanted = stack[sp - 1u];
                let count = rule_counts[count_index(rule_id, inst.operand)];
                var ok = 0u;
                for (var i = 0u; i < min(count, params.z); i = i + 1u) {{
                    if (rule_positions[pos_index(rule_id, inst.operand, i)] == wanted) {{ ok = 1u; }}
                }}
                stack[sp - 1u] = ok;
            }}
            case 26u: {{
                let hi = stack[sp - 1u];
                let lo = stack[sp - 2u];
                let count = rule_counts[count_index(rule_id, inst.operand)];
                var ok = 0u;
                for (var i = 0u; i < min(count, params.z); i = i + 1u) {{
                    let pos = rule_positions[pos_index(rule_id, inst.operand, i)];
                    if (pos >= lo && pos <= hi) {{ ok = 1u; }}
                }}
                sp = sp - 1u;
                stack[sp - 1u] = ok;
            }}
            case 27u: {{
                let body_len = inst.operand >> 16u;
                let range_upper = stack[sp - 1u];
                let range_lower = stack[sp - 2u];
                sp = sp - 2u;
                if (body_len > 0u && pc > 0xFFFFFFFFu - body_len) {{
                    return;
                }}
                let body_start = pc;
                pc = pc + body_len;
                var result = false;
                if (body_len > 0u && range_lower <= range_upper) {{
                    let body_end = body_start + body_len - 1u;
                    let iter_count = min(range_upper - range_lower, MAX_FOR_ITERATIONS - 1u) + 1u;
                    for (var iter = 0u; iter < iter_count; iter = iter + 1u) {{
                        let body_sp = eval_body(body_start, body_end, rule_id, sp, &stack);
                        if (body_sp == ABORT_SENTINEL) {{ return; }}
                        if (body_sp > sp && stack[body_sp - 1u] != 0u) {{
                            result = true;
                            break;
                        }}
                    }}
                }}
                if (sp >= STACK_SIZE) {{ return; }}
                stack[sp] = select(0u, 1u, result);
                sp = sp + 1u;
            }}
            case 28u: {{
                let body_len = inst.operand >> 16u;
                let range_upper = stack[sp - 1u];
                let range_lower = stack[sp - 2u];
                sp = sp - 2u;
                if (body_len > 0u && pc > 0xFFFFFFFFu - body_len) {{
                    return;
                }}
                let body_start = pc;
                pc = pc + body_len;
                var result = true;
                if (body_len > 0u && range_lower <= range_upper) {{
                    let body_end = body_start + body_len - 1u;
                    let iter_count = min(range_upper - range_lower, MAX_FOR_ITERATIONS - 1u) + 1u;
                    for (var iter = 0u; iter < iter_count; iter = iter + 1u) {{
                        let body_sp = eval_body(body_start, body_end, rule_id, sp, &stack);
                        if (body_sp == ABORT_SENTINEL) {{ return; }}
                        if (body_sp <= sp || stack[body_sp - 1u] == 0u) {{
                            result = false;
                            break;
                        }}
                    }}
                }}
                if (sp >= STACK_SIZE) {{ return; }}
                stack[sp] = select(0u, 1u, result);
                sp = sp + 1u;
            }}
            case 29u: {{}}
            case 40u: {{
                let n = inst.operand & 0xFFFFu;
                let body_len = inst.operand >> 16u;
                let range_upper = stack[sp - 1u];
                let range_lower = stack[sp - 2u];
                sp = sp - 2u;
                if (body_len > 0u && pc > 0xFFFFFFFFu - body_len) {{
                    return;
                }}
                let body_start = pc;
                pc = pc + body_len;
                var true_count = 0u;
                if (body_len > 0u && range_lower <= range_upper) {{
                    let body_end = body_start + body_len - 1u;
                    let iter_count = min(range_upper - range_lower, MAX_FOR_ITERATIONS - 1u) + 1u;
                    for (var iter = 0u; iter < iter_count; iter = iter + 1u) {{
                        let body_sp = eval_body(body_start, body_end, rule_id, sp, &stack);
                        if (body_sp == ABORT_SENTINEL) {{ return; }}
                        if (body_sp > sp && stack[body_sp - 1u] != 0u) {{
                            true_count = true_count + 1u;
                            if (true_count >= n) {{
                                break;
                            }}
                        }}
                    }}
                }}
                if (sp >= STACK_SIZE) {{ return; }}
                stack[sp] = select(0u, 1u, true_count >= n);
                sp = sp + 1u;
            }}
            case 41u: {{ if (sp >= STACK_SIZE) {{ return; }} stack[sp] = file_ctx.entropy_bucket; sp = sp + 1u; }}
            case 42u: {{ if (sp >= STACK_SIZE) {{ return; }} stack[sp] = file_ctx.is_pe; sp = sp + 1u; }}
            case 43u: {{ if (sp >= STACK_SIZE) {{ return; }} stack[sp] = file_ctx.is_dll; sp = sp + 1u; }}
            case 44u: {{ if (sp >= STACK_SIZE) {{ return; }} stack[sp] = file_ctx.num_sections; sp = sp + 1u; }}
            case 45u: {{ if (sp >= STACK_SIZE) {{ return; }} stack[sp] = file_ctx.num_imports; sp = sp + 1u; }}
            case 46u: {{ if (sp >= STACK_SIZE) {{ return; }} stack[sp] = file_ctx.entry_point_rva; sp = sp + 1u; }}
            case 47u: {{ if (sp >= STACK_SIZE) {{ return; }} stack[sp] = file_ctx.has_signature; sp = sp + 1u; }}
            case 48u: {{ if (sp >= STACK_SIZE) {{ return; }} stack[sp] = file_ctx.magic_u32; sp = sp + 1u; }}
            case 49u: {{ if (sp >= STACK_SIZE) {{ return; }} stack[sp] = file_ctx.is_64bit; sp = sp + 1u; }}
            case 30u: {{
                if (stack[sp - 1u] != 0u) {{
                    let idx = atomicAdd(&fired_results.count, 1u);
                    if (idx < 1024u) {{
                        fired_results.indices[idx] = rule_id;
                    }}
                }}
                return;
            }}
            case 39u: {{
                let offset = stack[sp - 1u];
                let size = inst.operand & 0xFFu;
                let is_signed = ((inst.operand >> 8u) & 1u) != 0u;
                let is_be = ((inst.operand >> 9u) & 1u) != 0u;
                
                var value = 0u;
                if (offset + size <= file_ctx.file_size) {{
                    let b0 = read_byte(offset);
                    if (size == 1u) {{
                        if (is_signed && (b0 & 0x80u) != 0u) {{
                            value = b0 | 0xFFFFFF00u;
                        }} else {{
                            value = b0;
                        }}
                    }} else if (size == 2u) {{
                        let b1 = read_byte(offset + 1u);
                        var v = 0u;
                        if (is_be) {{
                            v = (b0 << 8u) | b1;
                        }} else {{
                            v = (b1 << 8u) | b0;
                        }}
                        if (is_signed && (v & 0x8000u) != 0u) {{
                            value = v | 0xFFFF0000u;
                        }} else {{
                            value = v;
                        }}
                    }} else if (size == 4u) {{
                        let b1 = read_byte(offset + 1u);
                        let b2 = read_byte(offset + 2u);
                        let b3 = read_byte(offset + 3u);
                        if (is_be) {{
                            value = (b0 << 24u) | (b1 << 16u) | (b2 << 8u) | b3;
                        }} else {{
                            value = (b3 << 24u) | (b2 << 16u) | (b1 << 8u) | b0;
                        }}
                    }}
                }}
                stack[sp - 1u] = value;
            }}
            default: {{
                return;
            }}
        }}
    }}
    // Fell through without HALT — conservative false
}}
"#
    )
}