# oracle dump: pie::metal::pipeline::emit_readiness_msl
# source of truth: compiler/codegen/src/metal/ (the C++ oracle this was dumped from is deleted)
# emitter_version: 33
# @runtime: bytes=39611 fnv1a64=0x310a2f7d72f22fd5
# @grouped: bytes=1043 fnv1a64=0x429904221e8e963f
# @grouped was bytes=970 fnv1a64=0xa8ab9e1ebf8714d9 through emitter 41. The lane
#   record grew `attn_score_base` and `attn_score_row_stride` (+73 bytes of struct
#   text, no field before them moved), which is the grouped form's route to the
#   attention-score rectangle: that form binds no per-intrinsic buffer, so a
#   rectangle it must reach arrives as an address on the record. Unlike @runtime,
#   this prefix is not kernel-side text an `oracle-inputs/` copy could stand in
#   for -- `codegen::layout` declares the table and prints it, so the bytes are
#   the emitter's and the move belongs in the diff.
# cases: 45
# re-blessed: the C++ oracle (crates/engine-metal/csrc/src/pipeline/m1_codegen.cpp) is
#   deleted, so these bytes can no longer be re-derived from it. Three cases
#   below were re-blessed from Rust in the change that made the emitted corpus
#   compile: `crates/engine-metal/csrc/tests/msl_compile_test.cpp` put every kernel here
#   through Apple's shader compiler for the first time and found two rejections
#   the byte comparison could never see, because the oracle had them too.
#     * readiness/commit `wide_30` emitted `[[buffer(31)]]`, one past Metal's
#       highest buffer index; `METAL_M1_MAX_CHANNELS` now rejects it by name.
#     * three grouped fused regions passed the `const device bfloat*` logits
#       pointer into `ptir_m1_execute`'s `const device uchar*` parameter, which
#       MSL will not convert. They now cast.
#   Nothing else in this directory moved, and `msl_compile_test` is the standing
#   check that took over from the oracle for the question it could not answer.
#   re-blessed again (emitter 27): `m1_reduce_argmax` no longer materializes one
#   16-byte candidate per element into device `temporary` and tree-reduces it in
#   passes -- `m1_argmax_combine` is a strict total order, so a fold gives the
#   same answer -- and the bf16 logits intrinsic copies unrolled. Both loops run
#   on the single thread that owns a lane, so at a 248320 vocab they were ~250ms
#   per decode step against a ~4ms model forward. Only @runtime bytes moved.
#   re-blessed again (emitter 29): a generated region now gets a threadgroup per
#   lane instead of a thread, and the three things that walked a vocabulary-wide
#   row -- the logits gather, argmax, and the typed copy behind a reshape -- are
#   split across it. m1_copy_typed also moved from one device access per BYTE to
#   one per word. Between them a decode step's GPU time went 89ms -> 6.5ms.
#   re-blessed again (emitter 31): a same-footprint reshape is a view, so the
#   grouped emitter elides its copy and points the consumers at the source
#   value's offset. The runtime executed it as a byte-for-byte copy of the whole
#   value, which in the sampler's graph was a vocabulary-wide round trip through
#   device memory for nothing.
# REGENERATED: the body below is this compiler's output, not the recorded output of the source of truth above
#   reason: metal emitter 41->42, and the shared lane record is what moved. The GROUPED (M3) form binds no per-intrinsic buffer -- every rectangle it reads arrives as an address on the lane record -- so `attn_score` was the one id its table refused: the observability slab is the shell's own reservation (engine_metal::scores), not the readout arena, and no row displacement off lane.logits_base reaches it. LaneRecord now carries attn_score_base and attn_score_row_stride (appended, so every field before them keeps its offset; the record is 96 -> 112 bytes), emit_grouped_fused_region gathers the F32 planes from that base, and m3_intrinsic_bindable stops declining the id. What moved in these dumps is exactly the +73 bytes of struct text the two fields add to the M1 and M3 preambles, plus the emitter version spelled into ptir_m3_generic_{ready,commit}_v42. No emitted logic changed for any corpus case: the base corpus reads no score rectangle, so no kernel body here grew an arm. The bump is load-bearing rather than hygienic -- a library cached under 41 reads the 96-byte record, so lane 1 of any grouped table would start sixteen bytes early -- and CUDA bumps 24 -> 25 for the same reason on its own copy of the struct. The reason the record grew at all is the M2 ceiling: the score rectangle costs two of the twelve argument slots the channels grow into, so a ten-channel guest reading scores (trackb-h2o) was refused by a limit that had nothing to do with what it asked for.
=== empty
function_name: ptir_m1_empty_ready
channels: 0
ok: true
source: bytes=1574 prefix=none tail_bytes=1574

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_empty_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 0) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  status->state = 1;
}
=== end
=== flags_0
function_name: ptir_m1_flags_0_ready
channels: 1
channel[0]: requires_full=0 requires_empty=0 take=0 put=0 capacity=1
ok: true
source: bytes=2026 prefix=none tail_bytes=2026

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_flags_0_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 1) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
  }
  status->state = 1;
}
=== end
=== flags_1
function_name: ptir_m1_flags_1_ready
channels: 1
channel[0]: requires_full=1 requires_empty=0 take=0 put=0 capacity=1
ok: true
source: bytes=2101 prefix=none tail_bytes=2101

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_flags_1_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 1) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (tail <= head) { status->state = 2; status->fault = 1024; return; }
  }
  status->state = 1;
}
=== end
=== flags_2
function_name: ptir_m1_flags_2_ready
channels: 1
channel[0]: requires_full=0 requires_empty=1 take=0 put=0 capacity=1
ok: true
source: bytes=2107 prefix=none tail_bytes=2107

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_flags_2_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 1) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (tail - head >= 1ul) { status->state = 2; status->fault = 1152; return; }
  }
  status->state = 1;
}
=== end
=== flags_3
function_name: ptir_m1_flags_3_ready
channels: 1
channel[0]: requires_full=1 requires_empty=1 take=0 put=0 capacity=1
ok: true
source: bytes=2182 prefix=none tail_bytes=2182

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_flags_3_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 1) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (tail <= head) { status->state = 2; status->fault = 1024; return; }
    if (tail - head >= 1ul) { status->state = 2; status->fault = 1152; return; }
  }
  status->state = 1;
}
=== end
=== flags_4
function_name: ptir_m1_flags_4_ready
channels: 1
channel[0]: requires_full=0 requires_empty=0 take=1 put=0 capacity=1
ok: true
source: bytes=2026 prefix=none tail_bytes=2026

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_flags_4_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 1) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
  }
  status->state = 1;
}
=== end
=== flags_5
function_name: ptir_m1_flags_5_ready
channels: 1
channel[0]: requires_full=1 requires_empty=0 take=1 put=0 capacity=1
ok: true
source: bytes=2101 prefix=none tail_bytes=2101

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_flags_5_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 1) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (tail <= head) { status->state = 2; status->fault = 1024; return; }
  }
  status->state = 1;
}
=== end
=== flags_6
function_name: ptir_m1_flags_6_ready
channels: 1
channel[0]: requires_full=0 requires_empty=1 take=1 put=0 capacity=1
ok: true
source: bytes=2107 prefix=none tail_bytes=2107

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_flags_6_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 1) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (tail - head >= 1ul) { status->state = 2; status->fault = 1152; return; }
  }
  status->state = 1;
}
=== end
=== flags_7
function_name: ptir_m1_flags_7_ready
channels: 1
channel[0]: requires_full=1 requires_empty=1 take=1 put=0 capacity=1
ok: true
source: bytes=2182 prefix=none tail_bytes=2182

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_flags_7_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 1) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (tail <= head) { status->state = 2; status->fault = 1024; return; }
    if (tail - head >= 1ul) { status->state = 2; status->fault = 1152; return; }
  }
  status->state = 1;
}
=== end
=== flags_8
function_name: ptir_m1_flags_8_ready
channels: 1
channel[0]: requires_full=0 requires_empty=0 take=0 put=1 capacity=1
ok: true
source: bytes=2163 prefix=none tail_bytes=2163

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_flags_8_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 1) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 1ul + 0ul) { status->state = 2; status->fault = 1280; return; }
  }
  status->state = 1;
}
=== end
=== flags_9
function_name: ptir_m1_flags_9_ready
channels: 1
channel[0]: requires_full=1 requires_empty=0 take=0 put=1 capacity=1
ok: true
source: bytes=2238 prefix=none tail_bytes=2238

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_flags_9_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 1) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (tail <= head) { status->state = 2; status->fault = 1024; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 1ul + 0ul) { status->state = 2; status->fault = 1280; return; }
  }
  status->state = 1;
}
=== end
=== flags_10
function_name: ptir_m1_flags_10_ready
channels: 1
channel[0]: requires_full=0 requires_empty=1 take=0 put=1 capacity=1
ok: true
source: bytes=2245 prefix=none tail_bytes=2245

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_flags_10_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 1) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (tail - head >= 1ul) { status->state = 2; status->fault = 1152; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 1ul + 0ul) { status->state = 2; status->fault = 1280; return; }
  }
  status->state = 1;
}
=== end
=== flags_11
function_name: ptir_m1_flags_11_ready
channels: 1
channel[0]: requires_full=1 requires_empty=1 take=0 put=1 capacity=1
ok: true
source: bytes=2320 prefix=none tail_bytes=2320

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_flags_11_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 1) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (tail <= head) { status->state = 2; status->fault = 1024; return; }
    if (tail - head >= 1ul) { status->state = 2; status->fault = 1152; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 1ul + 0ul) { status->state = 2; status->fault = 1280; return; }
  }
  status->state = 1;
}
=== end
=== flags_12
function_name: ptir_m1_flags_12_ready
channels: 1
channel[0]: requires_full=0 requires_empty=0 take=1 put=1 capacity=1
ok: true
source: bytes=2164 prefix=none tail_bytes=2164

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_flags_12_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 1) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 1ul + 0ul) { status->state = 2; status->fault = 1280; return; }
  }
  status->state = 1;
}
=== end
=== flags_13
function_name: ptir_m1_flags_13_ready
channels: 1
channel[0]: requires_full=1 requires_empty=0 take=1 put=1 capacity=1
ok: true
source: bytes=2239 prefix=none tail_bytes=2239

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_flags_13_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 1) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (tail <= head) { status->state = 2; status->fault = 1024; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 1ul + 1ul) { status->state = 2; status->fault = 1280; return; }
  }
  status->state = 1;
}
=== end
=== flags_14
function_name: ptir_m1_flags_14_ready
channels: 1
channel[0]: requires_full=0 requires_empty=1 take=1 put=1 capacity=1
ok: true
source: bytes=2245 prefix=none tail_bytes=2245

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_flags_14_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 1) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (tail - head >= 1ul) { status->state = 2; status->fault = 1152; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 1ul + 0ul) { status->state = 2; status->fault = 1280; return; }
  }
  status->state = 1;
}
=== end
=== flags_15
function_name: ptir_m1_flags_15_ready
channels: 1
channel[0]: requires_full=1 requires_empty=1 take=1 put=1 capacity=1
ok: true
source: bytes=2320 prefix=none tail_bytes=2320

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_flags_15_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 1) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (tail <= head) { status->state = 2; status->fault = 1024; return; }
    if (tail - head >= 1ul) { status->state = 2; status->fault = 1152; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 1ul + 1ul) { status->state = 2; status->fault = 1280; return; }
  }
  status->state = 1;
}
=== end
=== capacity_0
function_name: ptir_m1_capacity_0_ready
channels: 1
channel[0]: requires_full=1 requires_empty=0 take=1 put=1 capacity=0
ok: true
source: bytes=2241 prefix=none tail_bytes=2241

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_capacity_0_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 1) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (tail <= head) { status->state = 2; status->fault = 1024; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 0ul + 1ul) { status->state = 2; status->fault = 1280; return; }
  }
  status->state = 1;
}
=== end
=== capacity_1
function_name: ptir_m1_capacity_1_ready
channels: 1
channel[0]: requires_full=1 requires_empty=0 take=1 put=1 capacity=1
ok: true
source: bytes=2241 prefix=none tail_bytes=2241

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_capacity_1_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 1) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (tail <= head) { status->state = 2; status->fault = 1024; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 1ul + 1ul) { status->state = 2; status->fault = 1280; return; }
  }
  status->state = 1;
}
=== end
=== capacity_2
function_name: ptir_m1_capacity_2_ready
channels: 1
channel[0]: requires_full=1 requires_empty=0 take=1 put=1 capacity=2
ok: true
source: bytes=2241 prefix=none tail_bytes=2241

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_capacity_2_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 1) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (tail <= head) { status->state = 2; status->fault = 1024; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 2ul + 1ul) { status->state = 2; status->fault = 1280; return; }
  }
  status->state = 1;
}
=== end
=== capacity_3
function_name: ptir_m1_capacity_3_ready
channels: 1
channel[0]: requires_full=1 requires_empty=0 take=1 put=1 capacity=3
ok: true
source: bytes=2241 prefix=none tail_bytes=2241

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_capacity_3_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 1) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (tail <= head) { status->state = 2; status->fault = 1024; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 3ul + 1ul) { status->state = 2; status->fault = 1280; return; }
  }
  status->state = 1;
}
=== end
=== capacity_7
function_name: ptir_m1_capacity_7_ready
channels: 1
channel[0]: requires_full=1 requires_empty=0 take=1 put=1 capacity=7
ok: true
source: bytes=2241 prefix=none tail_bytes=2241

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_capacity_7_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 1) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (tail <= head) { status->state = 2; status->fault = 1024; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 7ul + 1ul) { status->state = 2; status->fault = 1280; return; }
  }
  status->state = 1;
}
=== end
=== capacity_64
function_name: ptir_m1_capacity_64_ready
channels: 1
channel[0]: requires_full=1 requires_empty=0 take=1 put=1 capacity=64
ok: true
source: bytes=2243 prefix=none tail_bytes=2243

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_capacity_64_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 1) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (tail <= head) { status->state = 2; status->fault = 1024; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 64ul + 1ul) { status->state = 2; status->fault = 1280; return; }
  }
  status->state = 1;
}
=== end
=== capacity_65535
function_name: ptir_m1_capacity_65535_ready
channels: 1
channel[0]: requires_full=1 requires_empty=0 take=1 put=1 capacity=65535
ok: true
source: bytes=2249 prefix=none tail_bytes=2249

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_capacity_65535_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 1) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (tail <= head) { status->state = 2; status->fault = 1024; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 65535ul + 1ul) { status->state = 2; status->fault = 1280; return; }
  }
  status->state = 1;
}
=== end
=== capacity_4294967295
function_name: ptir_m1_capacity_4294967295_ready
channels: 1
channel[0]: requires_full=1 requires_empty=0 take=1 put=1 capacity=4294967295
ok: true
source: bytes=2259 prefix=none tail_bytes=2259

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_capacity_4294967295_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 1) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (tail <= head) { status->state = 2; status->fault = 1024; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 4294967295ul + 1ul) { status->state = 2; status->fault = 1280; return; }
  }
  status->state = 1;
}
=== end
=== pair_0
function_name: ptir_m1_pair_0_ready
channels: 2
channel[0]: requires_full=0 requires_empty=0 take=0 put=0 capacity=1
channel[1]: requires_full=1 requires_empty=1 take=1 put=1 capacity=3
ok: true
source: bytes=2768 prefix=none tail_bytes=2768

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_pair_0_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], device ulong* words_1 [[buffer(3)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 2) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
  }
  {
    const ulong head = words_1[0];
    const ulong tail = words_1[1];
    if (words_1[2] != 0ul || words_1[3] != 0ul || tail < head) { status->state = 3; status->fault = 513; return; }
    const ulong expected_head = slots[1].expected_head;
    const ulong expected_tail = slots[1].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 769; return; }
    if (tail <= head) { status->state = 2; status->fault = 1025; return; }
    if (tail - head >= 3ul) { status->state = 2; status->fault = 1153; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 3ul + 1ul) { status->state = 2; status->fault = 1281; return; }
  }
  status->state = 1;
}
=== end
=== pair_1
function_name: ptir_m1_pair_1_ready
channels: 2
channel[0]: requires_full=1 requires_empty=0 take=0 put=0 capacity=1
channel[1]: requires_full=0 requires_empty=1 take=1 put=1 capacity=3
ok: true
source: bytes=2768 prefix=none tail_bytes=2768

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_pair_1_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], device ulong* words_1 [[buffer(3)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 2) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (tail <= head) { status->state = 2; status->fault = 1024; return; }
  }
  {
    const ulong head = words_1[0];
    const ulong tail = words_1[1];
    if (words_1[2] != 0ul || words_1[3] != 0ul || tail < head) { status->state = 3; status->fault = 513; return; }
    const ulong expected_head = slots[1].expected_head;
    const ulong expected_tail = slots[1].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 769; return; }
    if (tail - head >= 3ul) { status->state = 2; status->fault = 1153; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 3ul + 0ul) { status->state = 2; status->fault = 1281; return; }
  }
  status->state = 1;
}
=== end
=== pair_2
function_name: ptir_m1_pair_2_ready
channels: 2
channel[0]: requires_full=0 requires_empty=1 take=0 put=0 capacity=1
channel[1]: requires_full=1 requires_empty=0 take=1 put=1 capacity=3
ok: true
source: bytes=2768 prefix=none tail_bytes=2768

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_pair_2_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], device ulong* words_1 [[buffer(3)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 2) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (tail - head >= 1ul) { status->state = 2; status->fault = 1152; return; }
  }
  {
    const ulong head = words_1[0];
    const ulong tail = words_1[1];
    if (words_1[2] != 0ul || words_1[3] != 0ul || tail < head) { status->state = 3; status->fault = 513; return; }
    const ulong expected_head = slots[1].expected_head;
    const ulong expected_tail = slots[1].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 769; return; }
    if (tail <= head) { status->state = 2; status->fault = 1025; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 3ul + 1ul) { status->state = 2; status->fault = 1281; return; }
  }
  status->state = 1;
}
=== end
=== pair_3
function_name: ptir_m1_pair_3_ready
channels: 2
channel[0]: requires_full=1 requires_empty=1 take=0 put=0 capacity=1
channel[1]: requires_full=0 requires_empty=0 take=1 put=1 capacity=3
ok: true
source: bytes=2768 prefix=none tail_bytes=2768

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_pair_3_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], device ulong* words_1 [[buffer(3)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 2) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (tail <= head) { status->state = 2; status->fault = 1024; return; }
    if (tail - head >= 1ul) { status->state = 2; status->fault = 1152; return; }
  }
  {
    const ulong head = words_1[0];
    const ulong tail = words_1[1];
    if (words_1[2] != 0ul || words_1[3] != 0ul || tail < head) { status->state = 3; status->fault = 513; return; }
    const ulong expected_head = slots[1].expected_head;
    const ulong expected_tail = slots[1].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 769; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 3ul + 0ul) { status->state = 2; status->fault = 1281; return; }
  }
  status->state = 1;
}
=== end
=== pair_4
function_name: ptir_m1_pair_4_ready
channels: 2
channel[0]: requires_full=0 requires_empty=0 take=1 put=0 capacity=1
channel[1]: requires_full=1 requires_empty=1 take=0 put=1 capacity=3
ok: true
source: bytes=2768 prefix=none tail_bytes=2768

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_pair_4_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], device ulong* words_1 [[buffer(3)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 2) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
  }
  {
    const ulong head = words_1[0];
    const ulong tail = words_1[1];
    if (words_1[2] != 0ul || words_1[3] != 0ul || tail < head) { status->state = 3; status->fault = 513; return; }
    const ulong expected_head = slots[1].expected_head;
    const ulong expected_tail = slots[1].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 769; return; }
    if (tail <= head) { status->state = 2; status->fault = 1025; return; }
    if (tail - head >= 3ul) { status->state = 2; status->fault = 1153; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 3ul + 0ul) { status->state = 2; status->fault = 1281; return; }
  }
  status->state = 1;
}
=== end
=== pair_5
function_name: ptir_m1_pair_5_ready
channels: 2
channel[0]: requires_full=1 requires_empty=0 take=1 put=0 capacity=1
channel[1]: requires_full=0 requires_empty=1 take=0 put=1 capacity=3
ok: true
source: bytes=2768 prefix=none tail_bytes=2768

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_pair_5_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], device ulong* words_1 [[buffer(3)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 2) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (tail <= head) { status->state = 2; status->fault = 1024; return; }
  }
  {
    const ulong head = words_1[0];
    const ulong tail = words_1[1];
    if (words_1[2] != 0ul || words_1[3] != 0ul || tail < head) { status->state = 3; status->fault = 513; return; }
    const ulong expected_head = slots[1].expected_head;
    const ulong expected_tail = slots[1].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 769; return; }
    if (tail - head >= 3ul) { status->state = 2; status->fault = 1153; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 3ul + 0ul) { status->state = 2; status->fault = 1281; return; }
  }
  status->state = 1;
}
=== end
=== pair_6
function_name: ptir_m1_pair_6_ready
channels: 2
channel[0]: requires_full=0 requires_empty=1 take=1 put=0 capacity=1
channel[1]: requires_full=1 requires_empty=0 take=0 put=1 capacity=3
ok: true
source: bytes=2768 prefix=none tail_bytes=2768

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_pair_6_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], device ulong* words_1 [[buffer(3)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 2) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (tail - head >= 1ul) { status->state = 2; status->fault = 1152; return; }
  }
  {
    const ulong head = words_1[0];
    const ulong tail = words_1[1];
    if (words_1[2] != 0ul || words_1[3] != 0ul || tail < head) { status->state = 3; status->fault = 513; return; }
    const ulong expected_head = slots[1].expected_head;
    const ulong expected_tail = slots[1].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 769; return; }
    if (tail <= head) { status->state = 2; status->fault = 1025; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 3ul + 0ul) { status->state = 2; status->fault = 1281; return; }
  }
  status->state = 1;
}
=== end
=== pair_7
function_name: ptir_m1_pair_7_ready
channels: 2
channel[0]: requires_full=1 requires_empty=1 take=1 put=0 capacity=1
channel[1]: requires_full=0 requires_empty=0 take=0 put=1 capacity=3
ok: true
source: bytes=2768 prefix=none tail_bytes=2768

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_pair_7_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], device ulong* words_1 [[buffer(3)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 2) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (tail <= head) { status->state = 2; status->fault = 1024; return; }
    if (tail - head >= 1ul) { status->state = 2; status->fault = 1152; return; }
  }
  {
    const ulong head = words_1[0];
    const ulong tail = words_1[1];
    if (words_1[2] != 0ul || words_1[3] != 0ul || tail < head) { status->state = 3; status->fault = 513; return; }
    const ulong expected_head = slots[1].expected_head;
    const ulong expected_tail = slots[1].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 769; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 3ul + 0ul) { status->state = 2; status->fault = 1281; return; }
  }
  status->state = 1;
}
=== end
=== pair_8
function_name: ptir_m1_pair_8_ready
channels: 2
channel[0]: requires_full=0 requires_empty=0 take=0 put=1 capacity=1
channel[1]: requires_full=1 requires_empty=1 take=1 put=0 capacity=3
ok: true
source: bytes=2768 prefix=none tail_bytes=2768

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_pair_8_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], device ulong* words_1 [[buffer(3)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 2) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 1ul + 0ul) { status->state = 2; status->fault = 1280; return; }
  }
  {
    const ulong head = words_1[0];
    const ulong tail = words_1[1];
    if (words_1[2] != 0ul || words_1[3] != 0ul || tail < head) { status->state = 3; status->fault = 513; return; }
    const ulong expected_head = slots[1].expected_head;
    const ulong expected_tail = slots[1].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 769; return; }
    if (tail <= head) { status->state = 2; status->fault = 1025; return; }
    if (tail - head >= 3ul) { status->state = 2; status->fault = 1153; return; }
  }
  status->state = 1;
}
=== end
=== pair_9
function_name: ptir_m1_pair_9_ready
channels: 2
channel[0]: requires_full=1 requires_empty=0 take=0 put=1 capacity=1
channel[1]: requires_full=0 requires_empty=1 take=1 put=0 capacity=3
ok: true
source: bytes=2768 prefix=none tail_bytes=2768

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_pair_9_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], device ulong* words_1 [[buffer(3)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 2) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (tail <= head) { status->state = 2; status->fault = 1024; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 1ul + 0ul) { status->state = 2; status->fault = 1280; return; }
  }
  {
    const ulong head = words_1[0];
    const ulong tail = words_1[1];
    if (words_1[2] != 0ul || words_1[3] != 0ul || tail < head) { status->state = 3; status->fault = 513; return; }
    const ulong expected_head = slots[1].expected_head;
    const ulong expected_tail = slots[1].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 769; return; }
    if (tail - head >= 3ul) { status->state = 2; status->fault = 1153; return; }
  }
  status->state = 1;
}
=== end
=== pair_10
function_name: ptir_m1_pair_10_ready
channels: 2
channel[0]: requires_full=0 requires_empty=1 take=0 put=1 capacity=1
channel[1]: requires_full=1 requires_empty=0 take=1 put=0 capacity=3
ok: true
source: bytes=2769 prefix=none tail_bytes=2769

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_pair_10_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], device ulong* words_1 [[buffer(3)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 2) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (tail - head >= 1ul) { status->state = 2; status->fault = 1152; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 1ul + 0ul) { status->state = 2; status->fault = 1280; return; }
  }
  {
    const ulong head = words_1[0];
    const ulong tail = words_1[1];
    if (words_1[2] != 0ul || words_1[3] != 0ul || tail < head) { status->state = 3; status->fault = 513; return; }
    const ulong expected_head = slots[1].expected_head;
    const ulong expected_tail = slots[1].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 769; return; }
    if (tail <= head) { status->state = 2; status->fault = 1025; return; }
  }
  status->state = 1;
}
=== end
=== pair_11
function_name: ptir_m1_pair_11_ready
channels: 2
channel[0]: requires_full=1 requires_empty=1 take=0 put=1 capacity=1
channel[1]: requires_full=0 requires_empty=0 take=1 put=0 capacity=3
ok: true
source: bytes=2769 prefix=none tail_bytes=2769

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_pair_11_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], device ulong* words_1 [[buffer(3)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 2) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (tail <= head) { status->state = 2; status->fault = 1024; return; }
    if (tail - head >= 1ul) { status->state = 2; status->fault = 1152; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 1ul + 0ul) { status->state = 2; status->fault = 1280; return; }
  }
  {
    const ulong head = words_1[0];
    const ulong tail = words_1[1];
    if (words_1[2] != 0ul || words_1[3] != 0ul || tail < head) { status->state = 3; status->fault = 513; return; }
    const ulong expected_head = slots[1].expected_head;
    const ulong expected_tail = slots[1].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 769; return; }
  }
  status->state = 1;
}
=== end
=== pair_12
function_name: ptir_m1_pair_12_ready
channels: 2
channel[0]: requires_full=0 requires_empty=0 take=1 put=1 capacity=1
channel[1]: requires_full=1 requires_empty=1 take=0 put=0 capacity=3
ok: true
source: bytes=2769 prefix=none tail_bytes=2769

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_pair_12_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], device ulong* words_1 [[buffer(3)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 2) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 1ul + 0ul) { status->state = 2; status->fault = 1280; return; }
  }
  {
    const ulong head = words_1[0];
    const ulong tail = words_1[1];
    if (words_1[2] != 0ul || words_1[3] != 0ul || tail < head) { status->state = 3; status->fault = 513; return; }
    const ulong expected_head = slots[1].expected_head;
    const ulong expected_tail = slots[1].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 769; return; }
    if (tail <= head) { status->state = 2; status->fault = 1025; return; }
    if (tail - head >= 3ul) { status->state = 2; status->fault = 1153; return; }
  }
  status->state = 1;
}
=== end
=== pair_13
function_name: ptir_m1_pair_13_ready
channels: 2
channel[0]: requires_full=1 requires_empty=0 take=1 put=1 capacity=1
channel[1]: requires_full=0 requires_empty=1 take=0 put=0 capacity=3
ok: true
source: bytes=2769 prefix=none tail_bytes=2769

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_pair_13_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], device ulong* words_1 [[buffer(3)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 2) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (tail <= head) { status->state = 2; status->fault = 1024; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 1ul + 1ul) { status->state = 2; status->fault = 1280; return; }
  }
  {
    const ulong head = words_1[0];
    const ulong tail = words_1[1];
    if (words_1[2] != 0ul || words_1[3] != 0ul || tail < head) { status->state = 3; status->fault = 513; return; }
    const ulong expected_head = slots[1].expected_head;
    const ulong expected_tail = slots[1].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 769; return; }
    if (tail - head >= 3ul) { status->state = 2; status->fault = 1153; return; }
  }
  status->state = 1;
}
=== end
=== pair_14
function_name: ptir_m1_pair_14_ready
channels: 2
channel[0]: requires_full=0 requires_empty=1 take=1 put=1 capacity=1
channel[1]: requires_full=1 requires_empty=0 take=0 put=0 capacity=3
ok: true
source: bytes=2769 prefix=none tail_bytes=2769

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_pair_14_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], device ulong* words_1 [[buffer(3)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 2) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (tail - head >= 1ul) { status->state = 2; status->fault = 1152; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 1ul + 0ul) { status->state = 2; status->fault = 1280; return; }
  }
  {
    const ulong head = words_1[0];
    const ulong tail = words_1[1];
    if (words_1[2] != 0ul || words_1[3] != 0ul || tail < head) { status->state = 3; status->fault = 513; return; }
    const ulong expected_head = slots[1].expected_head;
    const ulong expected_tail = slots[1].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 769; return; }
    if (tail <= head) { status->state = 2; status->fault = 1025; return; }
  }
  status->state = 1;
}
=== end
=== pair_15
function_name: ptir_m1_pair_15_ready
channels: 2
channel[0]: requires_full=1 requires_empty=1 take=1 put=1 capacity=1
channel[1]: requires_full=0 requires_empty=0 take=0 put=0 capacity=3
ok: true
source: bytes=2769 prefix=none tail_bytes=2769

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_pair_15_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], device ulong* words_1 [[buffer(3)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 2) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (tail <= head) { status->state = 2; status->fault = 1024; return; }
    if (tail - head >= 1ul) { status->state = 2; status->fault = 1152; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 1ul + 1ul) { status->state = 2; status->fault = 1280; return; }
  }
  {
    const ulong head = words_1[0];
    const ulong tail = words_1[1];
    if (words_1[2] != 0ul || words_1[3] != 0ul || tail < head) { status->state = 3; status->fault = 513; return; }
    const ulong expected_head = slots[1].expected_head;
    const ulong expected_tail = slots[1].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 769; return; }
  }
  status->state = 1;
}
=== end
=== mix5
function_name: ptir_m1_mix5_ready
channels: 5
channel[0]: requires_full=1 requires_empty=0 take=0 put=0 capacity=1
channel[1]: requires_full=0 requires_empty=0 take=0 put=1 capacity=2
channel[2]: requires_full=0 requires_empty=0 take=1 put=1 capacity=4
channel[3]: requires_full=0 requires_empty=1 take=0 put=0 capacity=8
channel[4]: requires_full=1 requires_empty=1 take=1 put=1 capacity=16
ok: true
source: bytes=4548 prefix=none tail_bytes=4548

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_mix5_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], device ulong* words_1 [[buffer(3)]], device ulong* words_2 [[buffer(4)]], device ulong* words_3 [[buffer(5)]], device ulong* words_4 [[buffer(6)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 5) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (tail <= head) { status->state = 2; status->fault = 1024; return; }
  }
  {
    const ulong head = words_1[0];
    const ulong tail = words_1[1];
    if (words_1[2] != 0ul || words_1[3] != 0ul || tail < head) { status->state = 3; status->fault = 513; return; }
    const ulong expected_head = slots[1].expected_head;
    const ulong expected_tail = slots[1].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 769; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 2ul + 0ul) { status->state = 2; status->fault = 1281; return; }
  }
  {
    const ulong head = words_2[0];
    const ulong tail = words_2[1];
    if (words_2[2] != 0ul || words_2[3] != 0ul || tail < head) { status->state = 3; status->fault = 514; return; }
    const ulong expected_head = slots[2].expected_head;
    const ulong expected_tail = slots[2].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 770; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 4ul + 0ul) { status->state = 2; status->fault = 1282; return; }
  }
  {
    const ulong head = words_3[0];
    const ulong tail = words_3[1];
    if (words_3[2] != 0ul || words_3[3] != 0ul || tail < head) { status->state = 3; status->fault = 515; return; }
    const ulong expected_head = slots[3].expected_head;
    const ulong expected_tail = slots[3].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 771; return; }
    if (tail - head >= 8ul) { status->state = 2; status->fault = 1155; return; }
  }
  {
    const ulong head = words_4[0];
    const ulong tail = words_4[1];
    if (words_4[2] != 0ul || words_4[3] != 0ul || tail < head) { status->state = 3; status->fault = 516; return; }
    const ulong expected_head = slots[4].expected_head;
    const ulong expected_tail = slots[4].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 772; return; }
    if (tail <= head) { status->state = 2; status->fault = 1028; return; }
    if (tail - head >= 16ul) { status->state = 2; status->fault = 1156; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 16ul + 1ul) { status->state = 2; status->fault = 1284; return; }
  }
  status->state = 1;
}
=== end
=== wide_28
function_name: ptir_m1_wide_28_ready
channels: 28
channel[0]: requires_full=1 requires_empty=0 take=0 put=0 capacity=1
channel[1]: requires_full=0 requires_empty=0 take=0 put=1 capacity=2
channel[2]: requires_full=1 requires_empty=1 take=1 put=1 capacity=3
channel[3]: requires_full=0 requires_empty=1 take=1 put=0 capacity=4
channel[4]: requires_full=1 requires_empty=0 take=1 put=1 capacity=5
channel[5]: requires_full=0 requires_empty=0 take=1 put=0 capacity=6
channel[6]: requires_full=1 requires_empty=1 take=0 put=1 capacity=7
channel[7]: requires_full=0 requires_empty=1 take=0 put=0 capacity=8
channel[8]: requires_full=1 requires_empty=0 take=0 put=1 capacity=9
channel[9]: requires_full=0 requires_empty=0 take=0 put=0 capacity=10
channel[10]: requires_full=1 requires_empty=1 take=1 put=0 capacity=11
channel[11]: requires_full=0 requires_empty=1 take=1 put=1 capacity=12
channel[12]: requires_full=1 requires_empty=0 take=1 put=0 capacity=13
channel[13]: requires_full=0 requires_empty=0 take=1 put=1 capacity=14
channel[14]: requires_full=1 requires_empty=1 take=0 put=0 capacity=15
channel[15]: requires_full=0 requires_empty=1 take=0 put=1 capacity=16
channel[16]: requires_full=1 requires_empty=0 take=0 put=0 capacity=17
channel[17]: requires_full=0 requires_empty=0 take=0 put=1 capacity=18
channel[18]: requires_full=1 requires_empty=1 take=1 put=1 capacity=19
channel[19]: requires_full=0 requires_empty=1 take=1 put=0 capacity=20
channel[20]: requires_full=1 requires_empty=0 take=1 put=1 capacity=21
channel[21]: requires_full=0 requires_empty=0 take=1 put=0 capacity=22
channel[22]: requires_full=1 requires_empty=1 take=0 put=1 capacity=23
channel[23]: requires_full=0 requires_empty=1 take=0 put=0 capacity=24
channel[24]: requires_full=1 requires_empty=0 take=0 put=1 capacity=25
channel[25]: requires_full=0 requires_empty=0 take=0 put=0 capacity=26
channel[26]: requires_full=1 requires_empty=1 take=1 put=0 capacity=27
channel[27]: requires_full=0 requires_empty=1 take=1 put=1 capacity=28
ok: true
source: bytes=18444 prefix=none tail_bytes=18444

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_wide_28_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], device ulong* words_1 [[buffer(3)]], device ulong* words_2 [[buffer(4)]], device ulong* words_3 [[buffer(5)]], device ulong* words_4 [[buffer(6)]], device ulong* words_5 [[buffer(7)]], device ulong* words_6 [[buffer(8)]], device ulong* words_7 [[buffer(9)]], device ulong* words_8 [[buffer(10)]], device ulong* words_9 [[buffer(11)]], device ulong* words_10 [[buffer(12)]], device ulong* words_11 [[buffer(13)]], device ulong* words_12 [[buffer(14)]], device ulong* words_13 [[buffer(15)]], device ulong* words_14 [[buffer(16)]], device ulong* words_15 [[buffer(17)]], device ulong* words_16 [[buffer(18)]], device ulong* words_17 [[buffer(19)]], device ulong* words_18 [[buffer(20)]], device ulong* words_19 [[buffer(21)]], device ulong* words_20 [[buffer(22)]], device ulong* words_21 [[buffer(23)]], device ulong* words_22 [[buffer(24)]], device ulong* words_23 [[buffer(25)]], device ulong* words_24 [[buffer(26)]], device ulong* words_25 [[buffer(27)]], device ulong* words_26 [[buffer(28)]], device ulong* words_27 [[buffer(29)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 28) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (tail <= head) { status->state = 2; status->fault = 1024; return; }
  }
  {
    const ulong head = words_1[0];
    const ulong tail = words_1[1];
    if (words_1[2] != 0ul || words_1[3] != 0ul || tail < head) { status->state = 3; status->fault = 513; return; }
    const ulong expected_head = slots[1].expected_head;
    const ulong expected_tail = slots[1].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 769; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 2ul + 0ul) { status->state = 2; status->fault = 1281; return; }
  }
  {
    const ulong head = words_2[0];
    const ulong tail = words_2[1];
    if (words_2[2] != 0ul || words_2[3] != 0ul || tail < head) { status->state = 3; status->fault = 514; return; }
    const ulong expected_head = slots[2].expected_head;
    const ulong expected_tail = slots[2].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 770; return; }
    if (tail <= head) { status->state = 2; status->fault = 1026; return; }
    if (tail - head >= 3ul) { status->state = 2; status->fault = 1154; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 3ul + 1ul) { status->state = 2; status->fault = 1282; return; }
  }
  {
    const ulong head = words_3[0];
    const ulong tail = words_3[1];
    if (words_3[2] != 0ul || words_3[3] != 0ul || tail < head) { status->state = 3; status->fault = 515; return; }
    const ulong expected_head = slots[3].expected_head;
    const ulong expected_tail = slots[3].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 771; return; }
    if (tail - head >= 4ul) { status->state = 2; status->fault = 1155; return; }
  }
  {
    const ulong head = words_4[0];
    const ulong tail = words_4[1];
    if (words_4[2] != 0ul || words_4[3] != 0ul || tail < head) { status->state = 3; status->fault = 516; return; }
    const ulong expected_head = slots[4].expected_head;
    const ulong expected_tail = slots[4].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 772; return; }
    if (tail <= head) { status->state = 2; status->fault = 1028; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 5ul + 1ul) { status->state = 2; status->fault = 1284; return; }
  }
  {
    const ulong head = words_5[0];
    const ulong tail = words_5[1];
    if (words_5[2] != 0ul || words_5[3] != 0ul || tail < head) { status->state = 3; status->fault = 517; return; }
    const ulong expected_head = slots[5].expected_head;
    const ulong expected_tail = slots[5].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 773; return; }
  }
  {
    const ulong head = words_6[0];
    const ulong tail = words_6[1];
    if (words_6[2] != 0ul || words_6[3] != 0ul || tail < head) { status->state = 3; status->fault = 518; return; }
    const ulong expected_head = slots[6].expected_head;
    const ulong expected_tail = slots[6].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 774; return; }
    if (tail <= head) { status->state = 2; status->fault = 1030; return; }
    if (tail - head >= 7ul) { status->state = 2; status->fault = 1158; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 7ul + 0ul) { status->state = 2; status->fault = 1286; return; }
  }
  {
    const ulong head = words_7[0];
    const ulong tail = words_7[1];
    if (words_7[2] != 0ul || words_7[3] != 0ul || tail < head) { status->state = 3; status->fault = 519; return; }
    const ulong expected_head = slots[7].expected_head;
    const ulong expected_tail = slots[7].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 775; return; }
    if (tail - head >= 8ul) { status->state = 2; status->fault = 1159; return; }
  }
  {
    const ulong head = words_8[0];
    const ulong tail = words_8[1];
    if (words_8[2] != 0ul || words_8[3] != 0ul || tail < head) { status->state = 3; status->fault = 520; return; }
    const ulong expected_head = slots[8].expected_head;
    const ulong expected_tail = slots[8].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 776; return; }
    if (tail <= head) { status->state = 2; status->fault = 1032; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 9ul + 0ul) { status->state = 2; status->fault = 1288; return; }
  }
  {
    const ulong head = words_9[0];
    const ulong tail = words_9[1];
    if (words_9[2] != 0ul || words_9[3] != 0ul || tail < head) { status->state = 3; status->fault = 521; return; }
    const ulong expected_head = slots[9].expected_head;
    const ulong expected_tail = slots[9].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 777; return; }
  }
  {
    const ulong head = words_10[0];
    const ulong tail = words_10[1];
    if (words_10[2] != 0ul || words_10[3] != 0ul || tail < head) { status->state = 3; status->fault = 522; return; }
    const ulong expected_head = slots[10].expected_head;
    const ulong expected_tail = slots[10].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 778; return; }
    if (tail <= head) { status->state = 2; status->fault = 1034; return; }
    if (tail - head >= 11ul) { status->state = 2; status->fault = 1162; return; }
  }
  {
    const ulong head = words_11[0];
    const ulong tail = words_11[1];
    if (words_11[2] != 0ul || words_11[3] != 0ul || tail < head) { status->state = 3; status->fault = 523; return; }
    const ulong expected_head = slots[11].expected_head;
    const ulong expected_tail = slots[11].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 779; return; }
    if (tail - head >= 12ul) { status->state = 2; status->fault = 1163; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 12ul + 0ul) { status->state = 2; status->fault = 1291; return; }
  }
  {
    const ulong head = words_12[0];
    const ulong tail = words_12[1];
    if (words_12[2] != 0ul || words_12[3] != 0ul || tail < head) { status->state = 3; status->fault = 524; return; }
    const ulong expected_head = slots[12].expected_head;
    const ulong expected_tail = slots[12].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 780; return; }
    if (tail <= head) { status->state = 2; status->fault = 1036; return; }
  }
  {
    const ulong head = words_13[0];
    const ulong tail = words_13[1];
    if (words_13[2] != 0ul || words_13[3] != 0ul || tail < head) { status->state = 3; status->fault = 525; return; }
    const ulong expected_head = slots[13].expected_head;
    const ulong expected_tail = slots[13].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 781; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 14ul + 0ul) { status->state = 2; status->fault = 1293; return; }
  }
  {
    const ulong head = words_14[0];
    const ulong tail = words_14[1];
    if (words_14[2] != 0ul || words_14[3] != 0ul || tail < head) { status->state = 3; status->fault = 526; return; }
    const ulong expected_head = slots[14].expected_head;
    const ulong expected_tail = slots[14].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 782; return; }
    if (tail <= head) { status->state = 2; status->fault = 1038; return; }
    if (tail - head >= 15ul) { status->state = 2; status->fault = 1166; return; }
  }
  {
    const ulong head = words_15[0];
    const ulong tail = words_15[1];
    if (words_15[2] != 0ul || words_15[3] != 0ul || tail < head) { status->state = 3; status->fault = 527; return; }
    const ulong expected_head = slots[15].expected_head;
    const ulong expected_tail = slots[15].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 783; return; }
    if (tail - head >= 16ul) { status->state = 2; status->fault = 1167; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 16ul + 0ul) { status->state = 2; status->fault = 1295; return; }
  }
  {
    const ulong head = words_16[0];
    const ulong tail = words_16[1];
    if (words_16[2] != 0ul || words_16[3] != 0ul || tail < head) { status->state = 3; status->fault = 528; return; }
    const ulong expected_head = slots[16].expected_head;
    const ulong expected_tail = slots[16].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 784; return; }
    if (tail <= head) { status->state = 2; status->fault = 1040; return; }
  }
  {
    const ulong head = words_17[0];
    const ulong tail = words_17[1];
    if (words_17[2] != 0ul || words_17[3] != 0ul || tail < head) { status->state = 3; status->fault = 529; return; }
    const ulong expected_head = slots[17].expected_head;
    const ulong expected_tail = slots[17].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 785; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 18ul + 0ul) { status->state = 2; status->fault = 1297; return; }
  }
  {
    const ulong head = words_18[0];
    const ulong tail = words_18[1];
    if (words_18[2] != 0ul || words_18[3] != 0ul || tail < head) { status->state = 3; status->fault = 530; return; }
    const ulong expected_head = slots[18].expected_head;
    const ulong expected_tail = slots[18].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 786; return; }
    if (tail <= head) { status->state = 2; status->fault = 1042; return; }
    if (tail - head >= 19ul) { status->state = 2; status->fault = 1170; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 19ul + 1ul) { status->state = 2; status->fault = 1298; return; }
  }
  {
    const ulong head = words_19[0];
    const ulong tail = words_19[1];
    if (words_19[2] != 0ul || words_19[3] != 0ul || tail < head) { status->state = 3; status->fault = 531; return; }
    const ulong expected_head = slots[19].expected_head;
    const ulong expected_tail = slots[19].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 787; return; }
    if (tail - head >= 20ul) { status->state = 2; status->fault = 1171; return; }
  }
  {
    const ulong head = words_20[0];
    const ulong tail = words_20[1];
    if (words_20[2] != 0ul || words_20[3] != 0ul || tail < head) { status->state = 3; status->fault = 532; return; }
    const ulong expected_head = slots[20].expected_head;
    const ulong expected_tail = slots[20].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 788; return; }
    if (tail <= head) { status->state = 2; status->fault = 1044; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 21ul + 1ul) { status->state = 2; status->fault = 1300; return; }
  }
  {
    const ulong head = words_21[0];
    const ulong tail = words_21[1];
    if (words_21[2] != 0ul || words_21[3] != 0ul || tail < head) { status->state = 3; status->fault = 533; return; }
    const ulong expected_head = slots[21].expected_head;
    const ulong expected_tail = slots[21].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 789; return; }
  }
  {
    const ulong head = words_22[0];
    const ulong tail = words_22[1];
    if (words_22[2] != 0ul || words_22[3] != 0ul || tail < head) { status->state = 3; status->fault = 534; return; }
    const ulong expected_head = slots[22].expected_head;
    const ulong expected_tail = slots[22].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 790; return; }
    if (tail <= head) { status->state = 2; status->fault = 1046; return; }
    if (tail - head >= 23ul) { status->state = 2; status->fault = 1174; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 23ul + 0ul) { status->state = 2; status->fault = 1302; return; }
  }
  {
    const ulong head = words_23[0];
    const ulong tail = words_23[1];
    if (words_23[2] != 0ul || words_23[3] != 0ul || tail < head) { status->state = 3; status->fault = 535; return; }
    const ulong expected_head = slots[23].expected_head;
    const ulong expected_tail = slots[23].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 791; return; }
    if (tail - head >= 24ul) { status->state = 2; status->fault = 1175; return; }
  }
  {
    const ulong head = words_24[0];
    const ulong tail = words_24[1];
    if (words_24[2] != 0ul || words_24[3] != 0ul || tail < head) { status->state = 3; status->fault = 536; return; }
    const ulong expected_head = slots[24].expected_head;
    const ulong expected_tail = slots[24].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 792; return; }
    if (tail <= head) { status->state = 2; status->fault = 1048; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 25ul + 0ul) { status->state = 2; status->fault = 1304; return; }
  }
  {
    const ulong head = words_25[0];
    const ulong tail = words_25[1];
    if (words_25[2] != 0ul || words_25[3] != 0ul || tail < head) { status->state = 3; status->fault = 537; return; }
    const ulong expected_head = slots[25].expected_head;
    const ulong expected_tail = slots[25].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 793; return; }
  }
  {
    const ulong head = words_26[0];
    const ulong tail = words_26[1];
    if (words_26[2] != 0ul || words_26[3] != 0ul || tail < head) { status->state = 3; status->fault = 538; return; }
    const ulong expected_head = slots[26].expected_head;
    const ulong expected_tail = slots[26].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 794; return; }
    if (tail <= head) { status->state = 2; status->fault = 1050; return; }
    if (tail - head >= 27ul) { status->state = 2; status->fault = 1178; return; }
  }
  {
    const ulong head = words_27[0];
    const ulong tail = words_27[1];
    if (words_27[2] != 0ul || words_27[3] != 0ul || tail < head) { status->state = 3; status->fault = 539; return; }
    const ulong expected_head = slots[27].expected_head;
    const ulong expected_tail = slots[27].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 795; return; }
    if (tail - head >= 28ul) { status->state = 2; status->fault = 1179; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 28ul + 0ul) { status->state = 2; status->fault = 1307; return; }
  }
  status->state = 1;
}
=== end
=== wide_29
function_name: ptir_m1_wide_29_ready
channels: 29
channel[0]: requires_full=1 requires_empty=0 take=0 put=0 capacity=1
channel[1]: requires_full=0 requires_empty=0 take=0 put=1 capacity=2
channel[2]: requires_full=1 requires_empty=1 take=1 put=1 capacity=3
channel[3]: requires_full=0 requires_empty=1 take=1 put=0 capacity=4
channel[4]: requires_full=1 requires_empty=0 take=1 put=1 capacity=5
channel[5]: requires_full=0 requires_empty=0 take=1 put=0 capacity=6
channel[6]: requires_full=1 requires_empty=1 take=0 put=1 capacity=7
channel[7]: requires_full=0 requires_empty=1 take=0 put=0 capacity=8
channel[8]: requires_full=1 requires_empty=0 take=0 put=1 capacity=9
channel[9]: requires_full=0 requires_empty=0 take=0 put=0 capacity=10
channel[10]: requires_full=1 requires_empty=1 take=1 put=0 capacity=11
channel[11]: requires_full=0 requires_empty=1 take=1 put=1 capacity=12
channel[12]: requires_full=1 requires_empty=0 take=1 put=0 capacity=13
channel[13]: requires_full=0 requires_empty=0 take=1 put=1 capacity=14
channel[14]: requires_full=1 requires_empty=1 take=0 put=0 capacity=15
channel[15]: requires_full=0 requires_empty=1 take=0 put=1 capacity=16
channel[16]: requires_full=1 requires_empty=0 take=0 put=0 capacity=17
channel[17]: requires_full=0 requires_empty=0 take=0 put=1 capacity=18
channel[18]: requires_full=1 requires_empty=1 take=1 put=1 capacity=19
channel[19]: requires_full=0 requires_empty=1 take=1 put=0 capacity=20
channel[20]: requires_full=1 requires_empty=0 take=1 put=1 capacity=21
channel[21]: requires_full=0 requires_empty=0 take=1 put=0 capacity=22
channel[22]: requires_full=1 requires_empty=1 take=0 put=1 capacity=23
channel[23]: requires_full=0 requires_empty=1 take=0 put=0 capacity=24
channel[24]: requires_full=1 requires_empty=0 take=0 put=1 capacity=25
channel[25]: requires_full=0 requires_empty=0 take=0 put=0 capacity=26
channel[26]: requires_full=1 requires_empty=1 take=1 put=0 capacity=27
channel[27]: requires_full=0 requires_empty=1 take=1 put=1 capacity=28
channel[28]: requires_full=1 requires_empty=0 take=1 put=0 capacity=29
ok: true
source: bytes=18977 prefix=none tail_bytes=18977

#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };
struct M1LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M1LaneRecord {
  ulong logits_base;
  uint logits_row_offset;
  uint logits_row_count;
  uint kv_len;
  uint page_count;
  uint row_count;
  uint token_count;
  uint sampled_rows;
  uint query_len;
  uint key_len;
  uint channel_slot_offset;
  ulong rng_state;
  ulong commit_slot;
  ulong active_row_mask;
  ulong sample_output_channel_mask;
  ulong row_valid;
  uint row_valid_offset;
  uint reserved0;
  ulong attn_score_base;
  uint attn_score_row_stride;
  uint reserved1;
};
struct M1LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
kernel void ptir_m1_wide_29_ready(device M1Status* status [[buffer(0)]], const device uchar* lane_bytes [[buffer(1)]], device ulong* words_0 [[buffer(2)]], device ulong* words_1 [[buffer(3)]], device ulong* words_2 [[buffer(4)]], device ulong* words_3 [[buffer(5)]], device ulong* words_4 [[buffer(6)]], device ulong* words_5 [[buffer(7)]], device ulong* words_6 [[buffer(8)]], device ulong* words_7 [[buffer(9)]], device ulong* words_8 [[buffer(10)]], device ulong* words_9 [[buffer(11)]], device ulong* words_10 [[buffer(12)]], device ulong* words_11 [[buffer(13)]], device ulong* words_12 [[buffer(14)]], device ulong* words_13 [[buffer(15)]], device ulong* words_14 [[buffer(16)]], device ulong* words_15 [[buffer(17)]], device ulong* words_16 [[buffer(18)]], device ulong* words_17 [[buffer(19)]], device ulong* words_18 [[buffer(20)]], device ulong* words_19 [[buffer(21)]], device ulong* words_20 [[buffer(22)]], device ulong* words_21 [[buffer(23)]], device ulong* words_22 [[buffer(24)]], device ulong* words_23 [[buffer(25)]], device ulong* words_24 [[buffer(26)]], device ulong* words_25 [[buffer(27)]], device ulong* words_26 [[buffer(28)]], device ulong* words_27 [[buffer(29)]], device ulong* words_28 [[buffer(30)]], uint gid [[thread_position_in_grid]]) {
  if (gid != 0) return;
  const device M1LaneHeader* header = reinterpret_cast<const device M1LaneHeader*>(lane_bytes);
  const device M1LaneChannelSlot* slots = reinterpret_cast<const device M1LaneChannelSlot*>(lane_bytes + sizeof(M1LaneHeader) + sizeof(M1LaneRecord));
  if (header->abi_version != 3 || header->lane_count != 1 || header->channel_count != 29) { status->state = 3; status->fault = 0x100; status->reserved0 = header->abi_version; status->reserved1 = (header->lane_count << 16) | header->channel_count; return; }
  status->state = 0; status->fault = 0;
  {
    const ulong head = words_0[0];
    const ulong tail = words_0[1];
    if (words_0[2] != 0ul || words_0[3] != 0ul || tail < head) { status->state = 3; status->fault = 512; return; }
    const ulong expected_head = slots[0].expected_head;
    const ulong expected_tail = slots[0].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 768; return; }
    if (tail <= head) { status->state = 2; status->fault = 1024; return; }
  }
  {
    const ulong head = words_1[0];
    const ulong tail = words_1[1];
    if (words_1[2] != 0ul || words_1[3] != 0ul || tail < head) { status->state = 3; status->fault = 513; return; }
    const ulong expected_head = slots[1].expected_head;
    const ulong expected_tail = slots[1].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 769; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 2ul + 0ul) { status->state = 2; status->fault = 1281; return; }
  }
  {
    const ulong head = words_2[0];
    const ulong tail = words_2[1];
    if (words_2[2] != 0ul || words_2[3] != 0ul || tail < head) { status->state = 3; status->fault = 514; return; }
    const ulong expected_head = slots[2].expected_head;
    const ulong expected_tail = slots[2].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 770; return; }
    if (tail <= head) { status->state = 2; status->fault = 1026; return; }
    if (tail - head >= 3ul) { status->state = 2; status->fault = 1154; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 3ul + 1ul) { status->state = 2; status->fault = 1282; return; }
  }
  {
    const ulong head = words_3[0];
    const ulong tail = words_3[1];
    if (words_3[2] != 0ul || words_3[3] != 0ul || tail < head) { status->state = 3; status->fault = 515; return; }
    const ulong expected_head = slots[3].expected_head;
    const ulong expected_tail = slots[3].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 771; return; }
    if (tail - head >= 4ul) { status->state = 2; status->fault = 1155; return; }
  }
  {
    const ulong head = words_4[0];
    const ulong tail = words_4[1];
    if (words_4[2] != 0ul || words_4[3] != 0ul || tail < head) { status->state = 3; status->fault = 516; return; }
    const ulong expected_head = slots[4].expected_head;
    const ulong expected_tail = slots[4].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 772; return; }
    if (tail <= head) { status->state = 2; status->fault = 1028; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 5ul + 1ul) { status->state = 2; status->fault = 1284; return; }
  }
  {
    const ulong head = words_5[0];
    const ulong tail = words_5[1];
    if (words_5[2] != 0ul || words_5[3] != 0ul || tail < head) { status->state = 3; status->fault = 517; return; }
    const ulong expected_head = slots[5].expected_head;
    const ulong expected_tail = slots[5].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 773; return; }
  }
  {
    const ulong head = words_6[0];
    const ulong tail = words_6[1];
    if (words_6[2] != 0ul || words_6[3] != 0ul || tail < head) { status->state = 3; status->fault = 518; return; }
    const ulong expected_head = slots[6].expected_head;
    const ulong expected_tail = slots[6].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 774; return; }
    if (tail <= head) { status->state = 2; status->fault = 1030; return; }
    if (tail - head >= 7ul) { status->state = 2; status->fault = 1158; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 7ul + 0ul) { status->state = 2; status->fault = 1286; return; }
  }
  {
    const ulong head = words_7[0];
    const ulong tail = words_7[1];
    if (words_7[2] != 0ul || words_7[3] != 0ul || tail < head) { status->state = 3; status->fault = 519; return; }
    const ulong expected_head = slots[7].expected_head;
    const ulong expected_tail = slots[7].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 775; return; }
    if (tail - head >= 8ul) { status->state = 2; status->fault = 1159; return; }
  }
  {
    const ulong head = words_8[0];
    const ulong tail = words_8[1];
    if (words_8[2] != 0ul || words_8[3] != 0ul || tail < head) { status->state = 3; status->fault = 520; return; }
    const ulong expected_head = slots[8].expected_head;
    const ulong expected_tail = slots[8].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 776; return; }
    if (tail <= head) { status->state = 2; status->fault = 1032; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 9ul + 0ul) { status->state = 2; status->fault = 1288; return; }
  }
  {
    const ulong head = words_9[0];
    const ulong tail = words_9[1];
    if (words_9[2] != 0ul || words_9[3] != 0ul || tail < head) { status->state = 3; status->fault = 521; return; }
    const ulong expected_head = slots[9].expected_head;
    const ulong expected_tail = slots[9].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 777; return; }
  }
  {
    const ulong head = words_10[0];
    const ulong tail = words_10[1];
    if (words_10[2] != 0ul || words_10[3] != 0ul || tail < head) { status->state = 3; status->fault = 522; return; }
    const ulong expected_head = slots[10].expected_head;
    const ulong expected_tail = slots[10].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 778; return; }
    if (tail <= head) { status->state = 2; status->fault = 1034; return; }
    if (tail - head >= 11ul) { status->state = 2; status->fault = 1162; return; }
  }
  {
    const ulong head = words_11[0];
    const ulong tail = words_11[1];
    if (words_11[2] != 0ul || words_11[3] != 0ul || tail < head) { status->state = 3; status->fault = 523; return; }
    const ulong expected_head = slots[11].expected_head;
    const ulong expected_tail = slots[11].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 779; return; }
    if (tail - head >= 12ul) { status->state = 2; status->fault = 1163; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 12ul + 0ul) { status->state = 2; status->fault = 1291; return; }
  }
  {
    const ulong head = words_12[0];
    const ulong tail = words_12[1];
    if (words_12[2] != 0ul || words_12[3] != 0ul || tail < head) { status->state = 3; status->fault = 524; return; }
    const ulong expected_head = slots[12].expected_head;
    const ulong expected_tail = slots[12].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 780; return; }
    if (tail <= head) { status->state = 2; status->fault = 1036; return; }
  }
  {
    const ulong head = words_13[0];
    const ulong tail = words_13[1];
    if (words_13[2] != 0ul || words_13[3] != 0ul || tail < head) { status->state = 3; status->fault = 525; return; }
    const ulong expected_head = slots[13].expected_head;
    const ulong expected_tail = slots[13].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 781; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 14ul + 0ul) { status->state = 2; status->fault = 1293; return; }
  }
  {
    const ulong head = words_14[0];
    const ulong tail = words_14[1];
    if (words_14[2] != 0ul || words_14[3] != 0ul || tail < head) { status->state = 3; status->fault = 526; return; }
    const ulong expected_head = slots[14].expected_head;
    const ulong expected_tail = slots[14].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 782; return; }
    if (tail <= head) { status->state = 2; status->fault = 1038; return; }
    if (tail - head >= 15ul) { status->state = 2; status->fault = 1166; return; }
  }
  {
    const ulong head = words_15[0];
    const ulong tail = words_15[1];
    if (words_15[2] != 0ul || words_15[3] != 0ul || tail < head) { status->state = 3; status->fault = 527; return; }
    const ulong expected_head = slots[15].expected_head;
    const ulong expected_tail = slots[15].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 783; return; }
    if (tail - head >= 16ul) { status->state = 2; status->fault = 1167; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 16ul + 0ul) { status->state = 2; status->fault = 1295; return; }
  }
  {
    const ulong head = words_16[0];
    const ulong tail = words_16[1];
    if (words_16[2] != 0ul || words_16[3] != 0ul || tail < head) { status->state = 3; status->fault = 528; return; }
    const ulong expected_head = slots[16].expected_head;
    const ulong expected_tail = slots[16].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 784; return; }
    if (tail <= head) { status->state = 2; status->fault = 1040; return; }
  }
  {
    const ulong head = words_17[0];
    const ulong tail = words_17[1];
    if (words_17[2] != 0ul || words_17[3] != 0ul || tail < head) { status->state = 3; status->fault = 529; return; }
    const ulong expected_head = slots[17].expected_head;
    const ulong expected_tail = slots[17].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 785; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 18ul + 0ul) { status->state = 2; status->fault = 1297; return; }
  }
  {
    const ulong head = words_18[0];
    const ulong tail = words_18[1];
    if (words_18[2] != 0ul || words_18[3] != 0ul || tail < head) { status->state = 3; status->fault = 530; return; }
    const ulong expected_head = slots[18].expected_head;
    const ulong expected_tail = slots[18].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 786; return; }
    if (tail <= head) { status->state = 2; status->fault = 1042; return; }
    if (tail - head >= 19ul) { status->state = 2; status->fault = 1170; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 19ul + 1ul) { status->state = 2; status->fault = 1298; return; }
  }
  {
    const ulong head = words_19[0];
    const ulong tail = words_19[1];
    if (words_19[2] != 0ul || words_19[3] != 0ul || tail < head) { status->state = 3; status->fault = 531; return; }
    const ulong expected_head = slots[19].expected_head;
    const ulong expected_tail = slots[19].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 787; return; }
    if (tail - head >= 20ul) { status->state = 2; status->fault = 1171; return; }
  }
  {
    const ulong head = words_20[0];
    const ulong tail = words_20[1];
    if (words_20[2] != 0ul || words_20[3] != 0ul || tail < head) { status->state = 3; status->fault = 532; return; }
    const ulong expected_head = slots[20].expected_head;
    const ulong expected_tail = slots[20].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 788; return; }
    if (tail <= head) { status->state = 2; status->fault = 1044; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 21ul + 1ul) { status->state = 2; status->fault = 1300; return; }
  }
  {
    const ulong head = words_21[0];
    const ulong tail = words_21[1];
    if (words_21[2] != 0ul || words_21[3] != 0ul || tail < head) { status->state = 3; status->fault = 533; return; }
    const ulong expected_head = slots[21].expected_head;
    const ulong expected_tail = slots[21].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 789; return; }
  }
  {
    const ulong head = words_22[0];
    const ulong tail = words_22[1];
    if (words_22[2] != 0ul || words_22[3] != 0ul || tail < head) { status->state = 3; status->fault = 534; return; }
    const ulong expected_head = slots[22].expected_head;
    const ulong expected_tail = slots[22].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 790; return; }
    if (tail <= head) { status->state = 2; status->fault = 1046; return; }
    if (tail - head >= 23ul) { status->state = 2; status->fault = 1174; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 23ul + 0ul) { status->state = 2; status->fault = 1302; return; }
  }
  {
    const ulong head = words_23[0];
    const ulong tail = words_23[1];
    if (words_23[2] != 0ul || words_23[3] != 0ul || tail < head) { status->state = 3; status->fault = 535; return; }
    const ulong expected_head = slots[23].expected_head;
    const ulong expected_tail = slots[23].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 791; return; }
    if (tail - head >= 24ul) { status->state = 2; status->fault = 1175; return; }
  }
  {
    const ulong head = words_24[0];
    const ulong tail = words_24[1];
    if (words_24[2] != 0ul || words_24[3] != 0ul || tail < head) { status->state = 3; status->fault = 536; return; }
    const ulong expected_head = slots[24].expected_head;
    const ulong expected_tail = slots[24].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 792; return; }
    if (tail <= head) { status->state = 2; status->fault = 1048; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 25ul + 0ul) { status->state = 2; status->fault = 1304; return; }
  }
  {
    const ulong head = words_25[0];
    const ulong tail = words_25[1];
    if (words_25[2] != 0ul || words_25[3] != 0ul || tail < head) { status->state = 3; status->fault = 537; return; }
    const ulong expected_head = slots[25].expected_head;
    const ulong expected_tail = slots[25].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 793; return; }
  }
  {
    const ulong head = words_26[0];
    const ulong tail = words_26[1];
    if (words_26[2] != 0ul || words_26[3] != 0ul || tail < head) { status->state = 3; status->fault = 538; return; }
    const ulong expected_head = slots[26].expected_head;
    const ulong expected_tail = slots[26].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 794; return; }
    if (tail <= head) { status->state = 2; status->fault = 1050; return; }
    if (tail - head >= 27ul) { status->state = 2; status->fault = 1178; return; }
  }
  {
    const ulong head = words_27[0];
    const ulong tail = words_27[1];
    if (words_27[2] != 0ul || words_27[3] != 0ul || tail < head) { status->state = 3; status->fault = 539; return; }
    const ulong expected_head = slots[27].expected_head;
    const ulong expected_tail = slots[27].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 795; return; }
    if (tail - head >= 28ul) { status->state = 2; status->fault = 1179; return; }
    if (expected_tail == ~0ul || tail != expected_tail || tail - head >= 28ul + 0ul) { status->state = 2; status->fault = 1307; return; }
  }
  {
    const ulong head = words_28[0];
    const ulong tail = words_28[1];
    if (words_28[2] != 0ul || words_28[3] != 0ul || tail < head) { status->state = 3; status->fault = 540; return; }
    const ulong expected_head = slots[28].expected_head;
    const ulong expected_tail = slots[28].expected_tail;
    if (expected_head != ~0ul && head != expected_head) { status->state = 2; status->fault = 796; return; }
    if (tail <= head) { status->state = 2; status->fault = 1052; return; }
  }
  status->state = 1;
}
=== end
=== wide_30
function_name: ptir_m1_wide_30_ready
channels: 30
channel[0]: requires_full=1 requires_empty=0 take=0 put=0 capacity=1
channel[1]: requires_full=0 requires_empty=0 take=0 put=1 capacity=2
channel[2]: requires_full=1 requires_empty=1 take=1 put=1 capacity=3
channel[3]: requires_full=0 requires_empty=1 take=1 put=0 capacity=4
channel[4]: requires_full=1 requires_empty=0 take=1 put=1 capacity=5
channel[5]: requires_full=0 requires_empty=0 take=1 put=0 capacity=6
channel[6]: requires_full=1 requires_empty=1 take=0 put=1 capacity=7
channel[7]: requires_full=0 requires_empty=1 take=0 put=0 capacity=8
channel[8]: requires_full=1 requires_empty=0 take=0 put=1 capacity=9
channel[9]: requires_full=0 requires_empty=0 take=0 put=0 capacity=10
channel[10]: requires_full=1 requires_empty=1 take=1 put=0 capacity=11
channel[11]: requires_full=0 requires_empty=1 take=1 put=1 capacity=12
channel[12]: requires_full=1 requires_empty=0 take=1 put=0 capacity=13
channel[13]: requires_full=0 requires_empty=0 take=1 put=1 capacity=14
channel[14]: requires_full=1 requires_empty=1 take=0 put=0 capacity=15
channel[15]: requires_full=0 requires_empty=1 take=0 put=1 capacity=16
channel[16]: requires_full=1 requires_empty=0 take=0 put=0 capacity=17
channel[17]: requires_full=0 requires_empty=0 take=0 put=1 capacity=18
channel[18]: requires_full=1 requires_empty=1 take=1 put=1 capacity=19
channel[19]: requires_full=0 requires_empty=1 take=1 put=0 capacity=20
channel[20]: requires_full=1 requires_empty=0 take=1 put=1 capacity=21
channel[21]: requires_full=0 requires_empty=0 take=1 put=0 capacity=22
channel[22]: requires_full=1 requires_empty=1 take=0 put=1 capacity=23
channel[23]: requires_full=0 requires_empty=1 take=0 put=0 capacity=24
channel[24]: requires_full=1 requires_empty=0 take=0 put=1 capacity=25
channel[25]: requires_full=0 requires_empty=0 take=0 put=0 capacity=26
channel[26]: requires_full=1 requires_empty=1 take=1 put=0 capacity=27
channel[27]: requires_full=0 requires_empty=1 take=1 put=1 capacity=28
channel[28]: requires_full=1 requires_empty=0 take=1 put=0 capacity=29
channel[29]: requires_full=0 requires_empty=0 take=1 put=1 capacity=30
ok: false
error: readiness kernel exceeds the 29-channel direct-binding limit
=== end
