pub fn shader_flow(max_stack: u32, simple_cases: &str) -> String {
let template = r#"
// 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: {
if (sp < 2u) { return ABORT_SENTINEL; }
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: {
if (sp < 2u) { return ABORT_SENTINEL; }
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: {
if (sp < 2u) { return ABORT_SENTINEL; }
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 27u: {
if (sp < 2u) { return; }
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: {
if (sp < 2u) { return; }
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 30u: {
if (sp == 0u) {
return;
}
if (stack[sp - 1u] != 0u) {
let idx = atomicAdd(&fired_results.count, 1u);
if (idx < 1024u) {
fired_results.indices[idx] = rule_id;
}
}
return;
}
case 40u: {
if (sp < 2u) { return; }
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;
}
__SIMPLE_CASES__
default: {
return;
}
}
}
}
"#;
template
.replace("__MAX_STACK__", &max_stack.to_string())
.replace("__SIMPLE_CASES__", simple_cases)
}