# oracle dump: pie::metal::pipeline::emit_grouped_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: 4
# 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.
=== ptir_m3_generic_ready_v42
function_name: ptir_m3_generic_ready_v42
ok: true
source: bytes=3050 prefix=none tail_bytes=3050
#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };

struct M3LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M3LaneRecord {
  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 M3LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
struct M3ChannelMeta {
  ulong words;
  uint capacity;
  uint flags;
};
struct M3GroupLayout {
  uint lane_count;
  uint value_count;
  uint scratch_stride;
  uint temporary_offset;
  uint vocab;
  uint reserved0;
  uint reserved1;
  uint reserved2;
};
struct M3RowMeta {
  uint offset;
  uint count;
  uint mtp_offset;
  uint reserved;
};
kernel void ptir_m3_generic_ready_v42(const device uchar* lane_bytes [[buffer(0)]], const device M3ChannelMeta* meta [[buffer(1)]], uint lane_index [[thread_position_in_grid]]) {
  const device M3LaneHeader* header = reinterpret_cast<const device M3LaneHeader*>(lane_bytes);
  if (lane_index >= header->lane_count) return;
  const device M3LaneRecord* lanes = reinterpret_cast<const device M3LaneRecord*>(lane_bytes + sizeof(M3LaneHeader));
  const device M3LaneChannelSlot* slots = reinterpret_cast<const device M3LaneChannelSlot*>(lane_bytes + sizeof(M3LaneHeader) + header->lane_count * sizeof(M3LaneRecord));
  const M3LaneRecord lane = lanes[lane_index];
  device M1Status* status = reinterpret_cast<device M1Status*>(lane.commit_slot);
  status->state = 0; status->fault = 0;
  for (uint channel = 0; channel < header->channel_count; ++channel) {
    const M3ChannelMeta m = meta[lane.channel_slot_offset + channel];
    if ((m.flags & 1u) == 0u) continue;
    const device ulong* words = reinterpret_cast<const device ulong*>(m.words);
    const ulong head = words[0], tail = words[1];
    const M3LaneChannelSlot slot = slots[lane.channel_slot_offset + channel];
    if (words[2] != 0ul || words[3] != 0ul || tail < head) { status->state = 3; status->fault = 0x700u + channel; return; }
    bool retry = slot.expected_head != ~0ul && head != slot.expected_head;
    retry = retry || (((m.flags & 2u) != 0u) && tail <= head);
    retry = retry || (((m.flags & 4u) != 0u) && tail - head >= ulong(m.capacity));
    if ((m.flags & 16u) != 0u) {
      const ulong credit = ((m.flags & 10u) == 10u) ? 1ul : 0ul;
      retry = retry || slot.expected_tail == ~0ul || tail != slot.expected_tail || tail - head >= ulong(m.capacity) + credit;
    }
    if (retry) { status->state = (m.flags & 32u) != 0u ? 3u : 2u; status->fault = 0x780u + channel; return; }
  }
  status->state = 1;
}
=== end
=== ptir_m3_generic_commit_v42
function_name: ptir_m3_generic_commit_v42
ok: true
source: bytes=3051 prefix=none tail_bytes=3051
#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };

struct M3LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M3LaneRecord {
  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 M3LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
struct M3ChannelMeta {
  ulong words;
  uint capacity;
  uint flags;
};
struct M3GroupLayout {
  uint lane_count;
  uint value_count;
  uint scratch_stride;
  uint temporary_offset;
  uint vocab;
  uint reserved0;
  uint reserved1;
  uint reserved2;
};
struct M3RowMeta {
  uint offset;
  uint count;
  uint mtp_offset;
  uint reserved;
};
kernel void ptir_m3_generic_commit_v42(const device uchar* lane_bytes [[buffer(0)]], const device M3ChannelMeta* meta [[buffer(1)]], uint lane_index [[thread_position_in_grid]]) {
  const device M3LaneHeader* header = reinterpret_cast<const device M3LaneHeader*>(lane_bytes);
  if (lane_index >= header->lane_count) return;
  const device M3LaneRecord* lanes = reinterpret_cast<const device M3LaneRecord*>(lane_bytes + sizeof(M3LaneHeader));
  const device M3LaneChannelSlot* slots = reinterpret_cast<const device M3LaneChannelSlot*>(lane_bytes + sizeof(M3LaneHeader) + header->lane_count * sizeof(M3LaneRecord));
  const M3LaneRecord lane = lanes[lane_index];
  device M1Status* status = reinterpret_cast<device M1Status*>(lane.commit_slot);
  status->state = 0; status->fault = 0;
  for (uint channel = 0; channel < header->channel_count; ++channel) {
    const M3ChannelMeta m = meta[lane.channel_slot_offset + channel];
    if ((m.flags & 1u) == 0u) continue;
    const device ulong* words = reinterpret_cast<const device ulong*>(m.words);
    const ulong head = words[0], tail = words[1];
    const M3LaneChannelSlot slot = slots[lane.channel_slot_offset + channel];
    if (words[2] != 0ul || words[3] != 0ul || tail < head) { status->state = 3; status->fault = 0x700u + channel; return; }
    bool retry = slot.expected_head != ~0ul && head != slot.expected_head;
    retry = retry || (((m.flags & 2u) != 0u) && tail <= head);
    retry = retry || (((m.flags & 4u) != 0u) && tail - head >= ulong(m.capacity));
    if ((m.flags & 16u) != 0u) {
      const ulong credit = ((m.flags & 10u) == 10u) ? 1ul : 0ul;
      retry = retry || slot.expected_tail == ~0ul || tail != slot.expected_tail || tail - head >= ulong(m.capacity) + credit;
    }
    if (retry) { status->state = (m.flags & 32u) != 0u ? 3u : 2u; status->fault = 0x780u + channel; return; }
  }
  status->state = 1;
}
=== end
=== k
function_name: k
ok: true
source: bytes=3026 prefix=none tail_bytes=3026
#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };

struct M3LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M3LaneRecord {
  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 M3LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
struct M3ChannelMeta {
  ulong words;
  uint capacity;
  uint flags;
};
struct M3GroupLayout {
  uint lane_count;
  uint value_count;
  uint scratch_stride;
  uint temporary_offset;
  uint vocab;
  uint reserved0;
  uint reserved1;
  uint reserved2;
};
struct M3RowMeta {
  uint offset;
  uint count;
  uint mtp_offset;
  uint reserved;
};
kernel void k(const device uchar* lane_bytes [[buffer(0)]], const device M3ChannelMeta* meta [[buffer(1)]], uint lane_index [[thread_position_in_grid]]) {
  const device M3LaneHeader* header = reinterpret_cast<const device M3LaneHeader*>(lane_bytes);
  if (lane_index >= header->lane_count) return;
  const device M3LaneRecord* lanes = reinterpret_cast<const device M3LaneRecord*>(lane_bytes + sizeof(M3LaneHeader));
  const device M3LaneChannelSlot* slots = reinterpret_cast<const device M3LaneChannelSlot*>(lane_bytes + sizeof(M3LaneHeader) + header->lane_count * sizeof(M3LaneRecord));
  const M3LaneRecord lane = lanes[lane_index];
  device M1Status* status = reinterpret_cast<device M1Status*>(lane.commit_slot);
  status->state = 0; status->fault = 0;
  for (uint channel = 0; channel < header->channel_count; ++channel) {
    const M3ChannelMeta m = meta[lane.channel_slot_offset + channel];
    if ((m.flags & 1u) == 0u) continue;
    const device ulong* words = reinterpret_cast<const device ulong*>(m.words);
    const ulong head = words[0], tail = words[1];
    const M3LaneChannelSlot slot = slots[lane.channel_slot_offset + channel];
    if (words[2] != 0ul || words[3] != 0ul || tail < head) { status->state = 3; status->fault = 0x700u + channel; return; }
    bool retry = slot.expected_head != ~0ul && head != slot.expected_head;
    retry = retry || (((m.flags & 2u) != 0u) && tail <= head);
    retry = retry || (((m.flags & 4u) != 0u) && tail - head >= ulong(m.capacity));
    if ((m.flags & 16u) != 0u) {
      const ulong credit = ((m.flags & 10u) == 10u) ? 1ul : 0ul;
      retry = retry || slot.expected_tail == ~0ul || tail != slot.expected_tail || tail - head >= ulong(m.capacity) + credit;
    }
    if (retry) { status->state = (m.flags & 32u) != 0u ? 3u : 2u; status->fault = 0x780u + channel; return; }
  }
  status->state = 1;
}
=== end
=== ptir_m3_0123456789abcdef_effects
function_name: ptir_m3_0123456789abcdef_effects
ok: true
source: bytes=3057 prefix=none tail_bytes=3057
#include <metal_stdlib>
using namespace metal;
struct M1Status { uint state; uint fault; uint reserved0; uint reserved1; };

struct M3LaneHeader { uint abi_version; uint lane_count; uint channel_count; uint flags; };
struct M3LaneRecord {
  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 M3LaneChannelSlot {
  ulong committed_cell;
  ulong pending_cell;
  ulong expected_head;
  ulong expected_tail;
};
struct M3ChannelMeta {
  ulong words;
  uint capacity;
  uint flags;
};
struct M3GroupLayout {
  uint lane_count;
  uint value_count;
  uint scratch_stride;
  uint temporary_offset;
  uint vocab;
  uint reserved0;
  uint reserved1;
  uint reserved2;
};
struct M3RowMeta {
  uint offset;
  uint count;
  uint mtp_offset;
  uint reserved;
};
kernel void ptir_m3_0123456789abcdef_effects(const device uchar* lane_bytes [[buffer(0)]], const device M3ChannelMeta* meta [[buffer(1)]], uint lane_index [[thread_position_in_grid]]) {
  const device M3LaneHeader* header = reinterpret_cast<const device M3LaneHeader*>(lane_bytes);
  if (lane_index >= header->lane_count) return;
  const device M3LaneRecord* lanes = reinterpret_cast<const device M3LaneRecord*>(lane_bytes + sizeof(M3LaneHeader));
  const device M3LaneChannelSlot* slots = reinterpret_cast<const device M3LaneChannelSlot*>(lane_bytes + sizeof(M3LaneHeader) + header->lane_count * sizeof(M3LaneRecord));
  const M3LaneRecord lane = lanes[lane_index];
  device M1Status* status = reinterpret_cast<device M1Status*>(lane.commit_slot);
  status->state = 0; status->fault = 0;
  for (uint channel = 0; channel < header->channel_count; ++channel) {
    const M3ChannelMeta m = meta[lane.channel_slot_offset + channel];
    if ((m.flags & 1u) == 0u) continue;
    const device ulong* words = reinterpret_cast<const device ulong*>(m.words);
    const ulong head = words[0], tail = words[1];
    const M3LaneChannelSlot slot = slots[lane.channel_slot_offset + channel];
    if (words[2] != 0ul || words[3] != 0ul || tail < head) { status->state = 3; status->fault = 0x700u + channel; return; }
    bool retry = slot.expected_head != ~0ul && head != slot.expected_head;
    retry = retry || (((m.flags & 2u) != 0u) && tail <= head);
    retry = retry || (((m.flags & 4u) != 0u) && tail - head >= ulong(m.capacity));
    if ((m.flags & 16u) != 0u) {
      const ulong credit = ((m.flags & 10u) == 10u) ? 1ul : 0ul;
      retry = retry || slot.expected_tail == ~0ul || tail != slot.expected_tail || tail - head >= ulong(m.capacity) + credit;
    }
    if (retry) { status->state = (m.flags & 32u) != 0u ? 3u : 2u; status->fault = 0x780u + channel; return; }
  }
  status->state = 1;
}
=== end
