# oracle dump: pie::metal::pipeline::emit_fused_region_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: 192
#   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: the M2 slot table: emit_fused_region now gives every intrinsic other than the trunk's logits an argument index of its own (down from buffer 30), so a region reading mtp_logits or mtp_drafts grows one buffer parameter and routes that op's a0 to it. The oracle emitted one buffer for all of them, which read the draft column as the trunk's logits; only the three mtp-reading cases move.
=== beam_epilogue#0 singleton#0
function_name: ptir_m2_b4fadd8446eb58dd_r0
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: fused region exceeds the 12-channel direct-binding limit
=== end
=== beam_epilogue#0 singleton#1
function_name: ptir_m2_b4fadd8446eb58dd_r1
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: fused region exceeds the 12-channel direct-binding limit
=== end
=== beam_epilogue#0 singleton#2
function_name: ptir_m2_b4fadd8446eb58dd_r2
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: fused region exceeds the 12-channel direct-binding limit
=== end
=== beam_epilogue#0 singleton#3
function_name: ptir_m2_b4fadd8446eb58dd_r3
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: fused region exceeds the 12-channel direct-binding limit
=== end
=== beam_epilogue#0 singleton#4
function_name: ptir_m2_b4fadd8446eb58dd_r4
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: fused region exceeds the 12-channel direct-binding limit
=== end
=== beam_epilogue#0 singleton#5
function_name: ptir_m2_b4fadd8446eb58dd_r5
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: fused region exceeds the 12-channel direct-binding limit
=== end
=== beam_epilogue#0 singleton#6
function_name: ptir_m2_b4fadd8446eb58dd_r6
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: fused region exceeds the 12-channel direct-binding limit
=== end
=== beam_epilogue#0 singleton#7
function_name: ptir_m2_b4fadd8446eb58dd_r7
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: fused region exceeds the 12-channel direct-binding limit
=== end
=== beam_epilogue#0 singleton#10
function_name: ptir_m2_b4fadd8446eb58dd_r10
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: fused region exceeds the 12-channel direct-binding limit
=== end
=== beam_epilogue#0 singleton#12
function_name: ptir_m2_b4fadd8446eb58dd_r12
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: fused region exceeds the 12-channel direct-binding limit
=== end
=== beam_epilogue#0 singleton#14
function_name: ptir_m2_b4fadd8446eb58dd_r14
region: library=1 library_op=1 schedule=3 nodes=1 inputs=1 outputs=2 sinks=0
ok: false
error: fused region exceeds the 12-channel direct-binding limit
=== end
=== beam_epilogue#0 singleton#15
function_name: ptir_m2_b4fadd8446eb58dd_r15
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: fused region exceeds the 12-channel direct-binding limit
=== end
=== beam_epilogue#0 singleton#16
function_name: ptir_m2_b4fadd8446eb58dd_r16
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: fused region exceeds the 12-channel direct-binding limit
=== end
=== beam_epilogue#0 singleton#18
function_name: ptir_m2_b4fadd8446eb58dd_r18
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: fused region exceeds the 12-channel direct-binding limit
=== end
=== beam_epilogue#0 singleton#25
function_name: ptir_m2_b4fadd8446eb58dd_r25
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: fused region exceeds the 12-channel direct-binding limit
=== end
=== beam_epilogue#0 singleton#26
function_name: ptir_m2_b4fadd8446eb58dd_r26
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: fused region exceeds the 12-channel direct-binding limit
=== end
=== beam_epilogue#0 singleton#28
function_name: ptir_m2_b4fadd8446eb58dd_r28
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: fused region exceeds the 12-channel direct-binding limit
=== end
=== beam_epilogue#0 singleton#30
function_name: ptir_m2_b4fadd8446eb58dd_r30
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: fused region exceeds the 12-channel direct-binding limit
=== end
=== beam_epilogue#0 singleton#31
function_name: ptir_m2_b4fadd8446eb58dd_r31
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: fused region exceeds the 12-channel direct-binding limit
=== end
=== beam_epilogue#0 singleton#35
function_name: ptir_m2_b4fadd8446eb58dd_r35
region: library=0 library_op=0 schedule=1 nodes=1 inputs=3 outputs=1 sinks=0
ok: false
error: fused region exceeds the 12-channel direct-binding limit
=== end
=== beam_epilogue#0 singleton#43
function_name: ptir_m2_b4fadd8446eb58dd_r43
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: fused region exceeds the 12-channel direct-binding limit
=== end
=== beam_epilogue#0 singleton#48
function_name: ptir_m2_b4fadd8446eb58dd_r48
region: library=0 library_op=0 schedule=0 nodes=1 inputs=1 outputs=0 sinks=1
ok: false
error: fused region exceeds the 12-channel direct-binding limit
=== end
=== beam_epilogue#0 singleton#75
function_name: ptir_m2_b4fadd8446eb58dd_r75
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: fused region exceeds the 12-channel direct-binding limit
=== end
=== beam_epilogue#0 singleton#76
function_name: ptir_m2_b4fadd8446eb58dd_r76
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: fused region exceeds the 12-channel direct-binding limit
=== end
=== beam_epilogue#0 fused#0
function_name: ptir_m2_b4fadd8446eb58dd_r0
region: library=0 library_op=0 schedule=1 nodes=14 inputs=0 outputs=1 sinks=0
ok: false
error: fused region exceeds the 12-channel direct-binding limit
=== end
=== beam_epilogue#0 fused#1
function_name: ptir_m2_b4fadd8446eb58dd_r1
region: library=1 library_op=1 schedule=3 nodes=1 inputs=1 outputs=2 sinks=0
ok: false
error: fused region exceeds the 12-channel direct-binding limit
=== end
=== beam_epilogue#0 fused#2
function_name: ptir_m2_b4fadd8446eb58dd_r2
region: library=0 library_op=0 schedule=1 nodes=67 inputs=2 outputs=0 sinks=15
ok: false
error: fused region exceeds the 12-channel direct-binding limit
=== end
=== counter_pingpong#0 singleton#0
function_name: ptir_m2_f4e949004f50e90e_r0
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41405 prefix=@runtime tail_bytes=906
kernel void ptir_m2_f4e949004f50e90e_r0(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  ptir_m1_execute(144u, status, descriptors, params + 0, current_0, scratch, scratch, scratch + offsets[0], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== counter_pingpong#0 singleton#1
function_name: ptir_m2_f4e949004f50e90e_r1
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41403 prefix=@runtime tail_bytes=904
kernel void ptir_m2_f4e949004f50e90e_r1(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  ptir_m1_execute(129u, status, descriptors, params + 1, scratch, scratch, scratch, scratch + offsets[1], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== counter_pingpong#0 singleton#2
function_name: ptir_m2_f4e949004f50e90e_r2
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: true
source: bytes=41428 prefix=@runtime tail_bytes=929
kernel void ptir_m2_f4e949004f50e90e_r2(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  ptir_m1_execute(16u, status, descriptors, params + 2, scratch + offsets[0], scratch + offsets[1], scratch, scratch + offsets[2], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== counter_pingpong#0 singleton#3
function_name: ptir_m2_f4e949004f50e90e_r3
region: library=0 library_op=0 schedule=0 nodes=1 inputs=1 outputs=0 sinks=1
ok: true
source: bytes=41430 prefix=@runtime tail_bytes=931
kernel void ptir_m2_f4e949004f50e90e_r3(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  ptir_m1_execute(146u, status, descriptors, params + 3, scratch + offsets[2], scratch, scratch, pending_0, scratch, temporary);
  if (status->state != 1) return;
  current_0 = pending_0;
}
=== end
=== counter_pingpong#0 fused#0
function_name: ptir_m2_f4e949004f50e90e_r0
region: library=0 library_op=0 schedule=1 nodes=5 inputs=0 outputs=0 sinks=2
ok: true
source: bytes=42128 prefix=@runtime tail_bytes=1629
kernel void ptir_m2_f4e949004f50e90e_r0(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  ptir_m1_execute(144u, status, descriptors, params + 0, current_0, scratch, scratch, scratch + offsets[0], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(129u, status, descriptors, params + 1, scratch, scratch, scratch, scratch + offsets[1], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(16u, status, descriptors, params + 2, scratch + offsets[0], scratch + offsets[1], scratch, scratch + offsets[2], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(146u, status, descriptors, params + 3, scratch + offsets[2], scratch, scratch, pending_0, scratch, temporary);
  if (status->state != 1) return;
  current_0 = pending_0;
  ptir_m1_execute(146u, status, descriptors, params + 4, scratch + offsets[2], scratch, scratch, pending_1, scratch, temporary);
  if (status->state != 1) return;
  current_1 = pending_1;
}
=== end
=== dfa_ingraph#0 singleton#0
function_name: ptir_m2_4f4b856e5ef3e8b4_r0
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41688 prefix=@runtime tail_bytes=1189
kernel void ptir_m2_4f4b856e5ef3e8b4_r0(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(160u, status, descriptors, params + 0, logits, scratch, scratch, scratch + offsets[0], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== dfa_ingraph#0 singleton#1
function_name: ptir_m2_4f4b856e5ef3e8b4_r1
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41691 prefix=@runtime tail_bytes=1192
kernel void ptir_m2_4f4b856e5ef3e8b4_r1(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(145u, status, descriptors, params + 1, current_0, scratch, scratch, scratch + offsets[1], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== dfa_ingraph#0 singleton#3
function_name: ptir_m2_4f4b856e5ef3e8b4_r3
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41691 prefix=@runtime tail_bytes=1192
kernel void ptir_m2_4f4b856e5ef3e8b4_r3(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(144u, status, descriptors, params + 3, current_2, scratch, scratch, scratch + offsets[3], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== dfa_ingraph#0 singleton#4
function_name: ptir_m2_4f4b856e5ef3e8b4_r4
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: true
source: bytes=41714 prefix=@runtime tail_bytes=1215
kernel void ptir_m2_4f4b856e5ef3e8b4_r4(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(96u, status, descriptors, params + 4, scratch + offsets[1], scratch + offsets[3], scratch, scratch + offsets[4], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== dfa_ingraph#0 singleton#5
function_name: ptir_m2_4f4b856e5ef3e8b4_r5
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41689 prefix=@runtime tail_bytes=1190
kernel void ptir_m2_4f4b856e5ef3e8b4_r5(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(129u, status, descriptors, params + 5, scratch, scratch, scratch, scratch + offsets[5], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== dfa_ingraph#0 singleton#6
function_name: ptir_m2_4f4b856e5ef3e8b4_r6
region: library=0 library_op=0 schedule=1 nodes=1 inputs=3 outputs=1 sinks=0
ok: true
source: bytes=41727 prefix=@runtime tail_bytes=1228
kernel void ptir_m2_4f4b856e5ef3e8b4_r6(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(32u, status, descriptors, params + 6, scratch + offsets[4], scratch + offsets[0], scratch + offsets[5], scratch + offsets[6], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== dfa_ingraph#0 singleton#7
function_name: ptir_m2_4f4b856e5ef3e8b4_r7
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41701 prefix=@runtime tail_bytes=1202
kernel void ptir_m2_4f4b856e5ef3e8b4_r7(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(51u, status, descriptors, params + 7, scratch + offsets[6], scratch, scratch, scratch + offsets[7], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== dfa_ingraph#0 singleton#8
function_name: ptir_m2_4f4b856e5ef3e8b4_r8
region: library=0 library_op=0 schedule=0 nodes=1 inputs=1 outputs=0 sinks=1
ok: true
source: bytes=41716 prefix=@runtime tail_bytes=1217
kernel void ptir_m2_4f4b856e5ef3e8b4_r8(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(146u, status, descriptors, params + 8, scratch + offsets[7], scratch, scratch, pending_3, scratch, temporary);
  if (status->state != 1) return;
  current_3 = pending_3;
}
=== end
=== dfa_ingraph#0 singleton#10
function_name: ptir_m2_4f4b856e5ef3e8b4_r10
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: true
source: bytes=41716 prefix=@runtime tail_bytes=1217
kernel void ptir_m2_4f4b856e5ef3e8b4_r10(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(18u, status, descriptors, params + 10, scratch + offsets[3], scratch + offsets[8], scratch, scratch + offsets[9], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== dfa_ingraph#0 singleton#11
function_name: ptir_m2_4f4b856e5ef3e8b4_r11
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41703 prefix=@runtime tail_bytes=1204
kernel void ptir_m2_4f4b856e5ef3e8b4_r11(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(7u, status, descriptors, params + 11, scratch + offsets[7], scratch, scratch, scratch + offsets[10], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== dfa_ingraph#0 singleton#12
function_name: ptir_m2_4f4b856e5ef3e8b4_r12
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: true
source: bytes=41718 prefix=@runtime tail_bytes=1219
kernel void ptir_m2_4f4b856e5ef3e8b4_r12(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(16u, status, descriptors, params + 12, scratch + offsets[9], scratch + offsets[10], scratch, scratch + offsets[11], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== dfa_ingraph#0 fused#0
function_name: ptir_m2_4f4b856e5ef3e8b4_r0
region: library=0 library_op=0 schedule=1 nodes=15 inputs=0 outputs=0 sinks=2
ok: true
source: bytes=44174 prefix=@runtime tail_bytes=3675
kernel void ptir_m2_4f4b856e5ef3e8b4_r0(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(160u, status, descriptors, params + 0, logits, scratch, scratch, scratch + offsets[0], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(145u, status, descriptors, params + 1, current_0, scratch, scratch, scratch + offsets[1], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(145u, status, descriptors, params + 2, current_1, scratch, scratch, scratch + offsets[2], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(144u, status, descriptors, params + 3, current_2, scratch, scratch, scratch + offsets[3], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(96u, status, descriptors, params + 4, scratch + offsets[1], scratch + offsets[3], scratch, scratch + offsets[4], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(129u, status, descriptors, params + 5, scratch, scratch, scratch, scratch + offsets[5], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(32u, status, descriptors, params + 6, scratch + offsets[4], scratch + offsets[0], scratch + offsets[5], scratch + offsets[6], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(51u, status, descriptors, params + 7, scratch + offsets[6], scratch, scratch, scratch + offsets[7], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(146u, status, descriptors, params + 8, scratch + offsets[7], scratch, scratch, pending_3, scratch, temporary);
  if (status->state != 1) return;
  current_3 = pending_3;
  ptir_m1_execute(129u, status, descriptors, params + 9, scratch, scratch, scratch, scratch + offsets[8], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(18u, status, descriptors, params + 10, scratch + offsets[3], scratch + offsets[8], scratch, scratch + offsets[9], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(7u, status, descriptors, params + 11, scratch + offsets[7], scratch, scratch, scratch + offsets[10], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(16u, status, descriptors, params + 12, scratch + offsets[9], scratch + offsets[10], scratch, scratch + offsets[11], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(96u, status, descriptors, params + 13, scratch + offsets[2], scratch + offsets[11], scratch, scratch + offsets[12], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(146u, status, descriptors, params + 14, scratch + offsets[12], scratch, scratch, pending_2, scratch, temporary);
  if (status->state != 1) return;
  current_2 = pending_2;
}
=== end
=== extern_contrastive#0 singleton#0
function_name: ptir_m2_d01f407b22863878_r0
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41260 prefix=@runtime tail_bytes=761
kernel void ptir_m2_d01f407b22863878_r0(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  ptir_m1_execute(160u, status, descriptors, params + 0, logits, scratch, scratch, scratch + offsets[0], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== extern_contrastive#0 singleton#1
function_name: ptir_m2_d01f407b22863878_r1
region: library=0 library_op=0 schedule=0 nodes=1 inputs=1 outputs=0 sinks=1
ok: true
source: bytes=41288 prefix=@runtime tail_bytes=789
kernel void ptir_m2_d01f407b22863878_r1(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  ptir_m1_execute(146u, status, descriptors, params + 1, scratch + offsets[0], scratch, scratch, pending_0, scratch, temporary);
  if (status->state != 1) return;
  current_0 = pending_0;
}
=== end
=== extern_contrastive#0 fused#0
function_name: ptir_m2_d01f407b22863878_r0
region: library=0 library_op=0 schedule=1 nodes=2 inputs=0 outputs=0 sinks=1
ok: true
source: bytes=41448 prefix=@runtime tail_bytes=949
kernel void ptir_m2_d01f407b22863878_r0(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  ptir_m1_execute(160u, status, descriptors, params + 0, logits, scratch, scratch, scratch + offsets[0], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(146u, status, descriptors, params + 1, scratch + offsets[0], scratch, scratch, pending_0, scratch, temporary);
  if (status->state != 1) return;
  current_0 = pending_0;
}
=== end
=== greedy_argmax#0 singleton#0
function_name: ptir_m2_f2566cc0c1e0529a_r0
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41402 prefix=@runtime tail_bytes=903
kernel void ptir_m2_f2566cc0c1e0529a_r0(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  ptir_m1_execute(160u, status, descriptors, params + 0, logits, scratch, scratch, scratch + offsets[0], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== greedy_argmax#0 singleton#1
function_name: ptir_m2_f2566cc0c1e0529a_r1
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41415 prefix=@runtime tail_bytes=916
kernel void ptir_m2_f2566cc0c1e0529a_r1(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  ptir_m1_execute(57u, status, descriptors, params + 1, scratch + offsets[0], scratch, scratch, scratch + offsets[1], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== greedy_argmax#0 singleton#2
function_name: ptir_m2_f2566cc0c1e0529a_r2
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41415 prefix=@runtime tail_bytes=916
kernel void ptir_m2_f2566cc0c1e0529a_r2(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  ptir_m1_execute(51u, status, descriptors, params + 2, scratch + offsets[1], scratch, scratch, scratch + offsets[2], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== greedy_argmax#0 singleton#4
function_name: ptir_m2_f2566cc0c1e0529a_r4
region: library=0 library_op=0 schedule=0 nodes=1 inputs=1 outputs=0 sinks=1
ok: true
source: bytes=41430 prefix=@runtime tail_bytes=931
kernel void ptir_m2_f2566cc0c1e0529a_r4(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  ptir_m1_execute(146u, status, descriptors, params + 4, scratch + offsets[3], scratch, scratch, pending_0, scratch, temporary);
  if (status->state != 1) return;
  current_0 = pending_0;
}
=== end
=== greedy_argmax#0 fused#0
function_name: ptir_m2_f2566cc0c1e0529a_r0
region: library=0 library_op=0 schedule=1 nodes=6 inputs=0 outputs=0 sinks=2
ok: true
source: bytes=42297 prefix=@runtime tail_bytes=1798
kernel void ptir_m2_f2566cc0c1e0529a_r0(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  ptir_m1_execute(160u, status, descriptors, params + 0, logits, scratch, scratch, scratch + offsets[0], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(57u, status, descriptors, params + 1, scratch + offsets[0], scratch, scratch, scratch + offsets[1], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(51u, status, descriptors, params + 2, scratch + offsets[1], scratch, scratch, scratch + offsets[2], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(57u, status, descriptors, params + 3, scratch + offsets[2], scratch, scratch, scratch + offsets[3], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(146u, status, descriptors, params + 4, scratch + offsets[3], scratch, scratch, pending_0, scratch, temporary);
  if (status->state != 1) return;
  current_0 = pending_0;
  ptir_m1_execute(146u, status, descriptors, params + 5, scratch + offsets[3], scratch, scratch, pending_1, scratch, temporary);
  if (status->state != 1) return;
  current_1 = pending_1;
}
=== end
=== lora_prologue#0 singleton#0
function_name: ptir_m2_26849b68dee54237_r0
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41405 prefix=@runtime tail_bytes=906
kernel void ptir_m2_26849b68dee54237_r0(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  ptir_m1_execute(145u, status, descriptors, params + 0, current_0, scratch, scratch, scratch + offsets[0], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== lora_prologue#0 singleton#2
function_name: ptir_m2_26849b68dee54237_r2
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41403 prefix=@runtime tail_bytes=904
kernel void ptir_m2_26849b68dee54237_r2(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  ptir_m1_execute(129u, status, descriptors, params + 2, scratch, scratch, scratch, scratch + offsets[2], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== lora_prologue#0 singleton#3
function_name: ptir_m2_26849b68dee54237_r3
region: library=1 library_op=5 schedule=3 nodes=1 inputs=3 outputs=0 sinks=0
ok: true
source: bytes=41429 prefix=@runtime tail_bytes=930
kernel void ptir_m2_26849b68dee54237_r3(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  ptir_m1_execute(162u, status, descriptors, params + 3, scratch + offsets[0], scratch + offsets[1], scratch + offsets[2], scratch, scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== lora_prologue#0 fused#0
function_name: ptir_m2_26849b68dee54237_r0
region: library=0 library_op=0 schedule=1 nodes=3 inputs=0 outputs=3 sinks=0
ok: true
source: bytes=41729 prefix=@runtime tail_bytes=1230
kernel void ptir_m2_26849b68dee54237_r0(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  ptir_m1_execute(145u, status, descriptors, params + 0, current_0, scratch, scratch, scratch + offsets[0], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(145u, status, descriptors, params + 1, current_1, scratch, scratch, scratch + offsets[1], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(129u, status, descriptors, params + 2, scratch, scratch, scratch, scratch + offsets[2], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== lora_prologue#0 fused#1
function_name: ptir_m2_26849b68dee54237_r1
region: library=1 library_op=5 schedule=3 nodes=1 inputs=3 outputs=0 sinks=0
ok: true
source: bytes=41429 prefix=@runtime tail_bytes=930
kernel void ptir_m2_26849b68dee54237_r1(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  ptir_m1_execute(162u, status, descriptors, params + 3, scratch + offsets[0], scratch + offsets[1], scratch + offsets[2], scratch, scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== matrix_mask_apply_packed#0 singleton#0
function_name: ptir_m2_6612fe00a51d7edc_r0
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41260 prefix=@runtime tail_bytes=761
kernel void ptir_m2_6612fe00a51d7edc_r0(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  ptir_m1_execute(160u, status, descriptors, params + 0, logits, scratch, scratch, scratch + offsets[0], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== matrix_mask_apply_packed#0 singleton#1
function_name: ptir_m2_6612fe00a51d7edc_r1
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41261 prefix=@runtime tail_bytes=762
kernel void ptir_m2_6612fe00a51d7edc_r1(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  ptir_m1_execute(129u, status, descriptors, params + 1, scratch, scratch, scratch, scratch + offsets[1], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== matrix_mask_apply_packed#0 singleton#2
function_name: ptir_m2_6612fe00a51d7edc_r2
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41273 prefix=@runtime tail_bytes=774
kernel void ptir_m2_6612fe00a51d7edc_r2(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  ptir_m1_execute(57u, status, descriptors, params + 2, scratch + offsets[1], scratch, scratch, scratch + offsets[2], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== matrix_mask_apply_packed#0 singleton#3
function_name: ptir_m2_6612fe00a51d7edc_r3
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: true
source: bytes=41287 prefix=@runtime tail_bytes=788
kernel void ptir_m2_6612fe00a51d7edc_r3(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  ptir_m1_execute(101u, status, descriptors, params + 3, scratch + offsets[0], scratch + offsets[2], scratch, scratch + offsets[3], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== matrix_mask_apply_packed#0 singleton#4
function_name: ptir_m2_6612fe00a51d7edc_r4
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41273 prefix=@runtime tail_bytes=774
kernel void ptir_m2_6612fe00a51d7edc_r4(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  ptir_m1_execute(51u, status, descriptors, params + 4, scratch + offsets[3], scratch, scratch, scratch + offsets[4], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== matrix_mask_apply_packed#0 singleton#5
function_name: ptir_m2_6612fe00a51d7edc_r5
region: library=0 library_op=0 schedule=0 nodes=1 inputs=1 outputs=0 sinks=1
ok: true
source: bytes=41288 prefix=@runtime tail_bytes=789
kernel void ptir_m2_6612fe00a51d7edc_r5(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  ptir_m1_execute(146u, status, descriptors, params + 5, scratch + offsets[4], scratch, scratch, pending_0, scratch, temporary);
  if (status->state != 1) return;
  current_0 = pending_0;
}
=== end
=== matrix_mask_apply_packed#0 fused#0
function_name: ptir_m2_6612fe00a51d7edc_r0
region: library=0 library_op=0 schedule=1 nodes=6 inputs=0 outputs=0 sinks=1
ok: true
source: bytes=42142 prefix=@runtime tail_bytes=1643
kernel void ptir_m2_6612fe00a51d7edc_r0(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  ptir_m1_execute(160u, status, descriptors, params + 0, logits, scratch, scratch, scratch + offsets[0], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(129u, status, descriptors, params + 1, scratch, scratch, scratch, scratch + offsets[1], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(57u, status, descriptors, params + 2, scratch + offsets[1], scratch, scratch, scratch + offsets[2], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(101u, status, descriptors, params + 3, scratch + offsets[0], scratch + offsets[2], scratch, scratch + offsets[3], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(51u, status, descriptors, params + 4, scratch + offsets[3], scratch, scratch, scratch + offsets[4], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(146u, status, descriptors, params + 5, scratch + offsets[4], scratch, scratch, pending_0, scratch, temporary);
  if (status->state != 1) return;
  current_0 = pending_0;
}
=== end
=== matrix_select_mask#0 singleton#0
function_name: ptir_m2_101abcfb79b08b89_r0
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41402 prefix=@runtime tail_bytes=903
kernel void ptir_m2_101abcfb79b08b89_r0(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  ptir_m1_execute(160u, status, descriptors, params + 0, logits, scratch, scratch, scratch + offsets[0], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== matrix_select_mask#0 singleton#1
function_name: ptir_m2_101abcfb79b08b89_r1
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41405 prefix=@runtime tail_bytes=906
kernel void ptir_m2_101abcfb79b08b89_r1(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  ptir_m1_execute(144u, status, descriptors, params + 1, current_0, scratch, scratch, scratch + offsets[1], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== matrix_select_mask#0 singleton#2
function_name: ptir_m2_101abcfb79b08b89_r2
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41403 prefix=@runtime tail_bytes=904
kernel void ptir_m2_101abcfb79b08b89_r2(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  ptir_m1_execute(129u, status, descriptors, params + 2, scratch, scratch, scratch, scratch + offsets[2], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== matrix_select_mask#0 singleton#3
function_name: ptir_m2_101abcfb79b08b89_r3
region: library=0 library_op=0 schedule=1 nodes=1 inputs=3 outputs=1 sinks=0
ok: true
source: bytes=41441 prefix=@runtime tail_bytes=942
kernel void ptir_m2_101abcfb79b08b89_r3(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  ptir_m1_execute(32u, status, descriptors, params + 3, scratch + offsets[1], scratch + offsets[0], scratch + offsets[2], scratch + offsets[3], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== matrix_select_mask#0 singleton#4
function_name: ptir_m2_101abcfb79b08b89_r4
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41415 prefix=@runtime tail_bytes=916
kernel void ptir_m2_101abcfb79b08b89_r4(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  ptir_m1_execute(51u, status, descriptors, params + 4, scratch + offsets[3], scratch, scratch, scratch + offsets[4], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== matrix_select_mask#0 singleton#5
function_name: ptir_m2_101abcfb79b08b89_r5
region: library=0 library_op=0 schedule=0 nodes=1 inputs=1 outputs=0 sinks=1
ok: true
source: bytes=41430 prefix=@runtime tail_bytes=931
kernel void ptir_m2_101abcfb79b08b89_r5(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  ptir_m1_execute(146u, status, descriptors, params + 5, scratch + offsets[4], scratch, scratch, pending_1, scratch, temporary);
  if (status->state != 1) return;
  current_1 = pending_1;
}
=== end
=== matrix_select_mask#0 fused#0
function_name: ptir_m2_101abcfb79b08b89_r0
region: library=0 library_op=0 schedule=1 nodes=6 inputs=0 outputs=0 sinks=1
ok: true
source: bytes=42286 prefix=@runtime tail_bytes=1787
kernel void ptir_m2_101abcfb79b08b89_r0(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  ptir_m1_execute(160u, status, descriptors, params + 0, logits, scratch, scratch, scratch + offsets[0], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(144u, status, descriptors, params + 1, current_0, scratch, scratch, scratch + offsets[1], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(129u, status, descriptors, params + 2, scratch, scratch, scratch, scratch + offsets[2], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(32u, status, descriptors, params + 3, scratch + offsets[1], scratch + offsets[0], scratch + offsets[2], scratch + offsets[3], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(51u, status, descriptors, params + 4, scratch + offsets[3], scratch, scratch, scratch + offsets[4], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(146u, status, descriptors, params + 5, scratch + offsets[4], scratch, scratch, pending_1, scratch, temporary);
  if (status->state != 1) return;
  current_1 = pending_1;
}
=== end
=== mtp_verify_tail#0 singleton#0
function_name: ptir_m2_963137d4ffed0445_r0
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41688 prefix=@runtime tail_bytes=1189
kernel void ptir_m2_963137d4ffed0445_r0(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(160u, status, descriptors, params + 0, logits, scratch, scratch, scratch + offsets[0], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== mtp_verify_tail#0 singleton#2
function_name: ptir_m2_963137d4ffed0445_r2
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41691 prefix=@runtime tail_bytes=1192
kernel void ptir_m2_963137d4ffed0445_r2(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(144u, status, descriptors, params + 2, current_0, scratch, scratch, scratch + offsets[2], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== mtp_verify_tail#0 singleton#3
function_name: ptir_m2_963137d4ffed0445_r3
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41689 prefix=@runtime tail_bytes=1190
kernel void ptir_m2_963137d4ffed0445_r3(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(129u, status, descriptors, params + 3, scratch, scratch, scratch, scratch + offsets[3], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== mtp_verify_tail#0 singleton#4
function_name: ptir_m2_963137d4ffed0445_r4
region: library=0 library_op=0 schedule=1 nodes=1 inputs=3 outputs=1 sinks=0
ok: true
source: bytes=41727 prefix=@runtime tail_bytes=1228
kernel void ptir_m2_963137d4ffed0445_r4(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(32u, status, descriptors, params + 4, scratch + offsets[2], scratch + offsets[0], scratch + offsets[3], scratch + offsets[4], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== mtp_verify_tail#0 singleton#5
function_name: ptir_m2_963137d4ffed0445_r5
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41701 prefix=@runtime tail_bytes=1202
kernel void ptir_m2_963137d4ffed0445_r5(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(51u, status, descriptors, params + 5, scratch + offsets[4], scratch, scratch, scratch + offsets[5], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== mtp_verify_tail#0 singleton#7
function_name: ptir_m2_963137d4ffed0445_r7
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41689 prefix=@runtime tail_bytes=1190
kernel void ptir_m2_963137d4ffed0445_r7(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(100u, status, descriptors, params + 7, scratch, scratch, scratch, scratch + offsets[7], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== mtp_verify_tail#0 singleton#8
function_name: ptir_m2_963137d4ffed0445_r8
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: true
source: bytes=41714 prefix=@runtime tail_bytes=1215
kernel void ptir_m2_963137d4ffed0445_r8(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(96u, status, descriptors, params + 8, scratch + offsets[5], scratch + offsets[7], scratch, scratch + offsets[8], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== mtp_verify_tail#0 singleton#9
function_name: ptir_m2_963137d4ffed0445_r9
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: true
source: bytes=41714 prefix=@runtime tail_bytes=1215
kernel void ptir_m2_963137d4ffed0445_r9(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(24u, status, descriptors, params + 9, scratch + offsets[6], scratch + offsets[8], scratch, scratch + offsets[9], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== mtp_verify_tail#0 singleton#13
function_name: ptir_m2_963137d4ffed0445_r13
region: library=1 library_op=3 schedule=3 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41705 prefix=@runtime tail_bytes=1206
kernel void ptir_m2_963137d4ffed0445_r13(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(65u, status, descriptors, params + 13, scratch + offsets[12], scratch, scratch, scratch + offsets[13], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== mtp_verify_tail#0 singleton#14
function_name: ptir_m2_963137d4ffed0445_r14
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41705 prefix=@runtime tail_bytes=1206
kernel void ptir_m2_963137d4ffed0445_r14(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(48u, status, descriptors, params + 14, scratch + offsets[13], scratch, scratch, scratch + offsets[14], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== mtp_verify_tail#0 singleton#15
function_name: ptir_m2_963137d4ffed0445_r15
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41704 prefix=@runtime tail_bytes=1205
kernel void ptir_m2_963137d4ffed0445_r15(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(7u, status, descriptors, params + 15, scratch + offsets[14], scratch, scratch, scratch + offsets[15], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== mtp_verify_tail#0 singleton#16
function_name: ptir_m2_963137d4ffed0445_r16
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41705 prefix=@runtime tail_bytes=1206
kernel void ptir_m2_963137d4ffed0445_r16(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(56u, status, descriptors, params + 16, scratch + offsets[15], scratch, scratch, scratch + offsets[16], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== mtp_verify_tail#0 singleton#18
function_name: ptir_m2_963137d4ffed0445_r18
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: true
source: bytes=41719 prefix=@runtime tail_bytes=1220
kernel void ptir_m2_963137d4ffed0445_r18(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(23u, status, descriptors, params + 18, scratch + offsets[16], scratch + offsets[17], scratch, scratch + offsets[18], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== mtp_verify_tail#0 singleton#21
function_name: ptir_m2_963137d4ffed0445_r21
region: library=0 library_op=0 schedule=0 nodes=1 inputs=1 outputs=0 sinks=1
ok: true
source: bytes=41719 prefix=@runtime tail_bytes=1220
kernel void ptir_m2_963137d4ffed0445_r21(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(146u, status, descriptors, params + 21, scratch + offsets[20], scratch, scratch, pending_2, scratch, temporary);
  if (status->state != 1) return;
  current_2 = pending_2;
}
=== end
=== mtp_verify_tail#0 fused#0
function_name: ptir_m2_963137d4ffed0445_r0
region: library=0 library_op=0 schedule=1 nodes=13 inputs=0 outputs=3 sinks=0
ok: true
source: bytes=43826 prefix=@runtime tail_bytes=3327
kernel void ptir_m2_963137d4ffed0445_r0(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* intrinsic_1 [[buffer(30)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(160u, status, descriptors, params + 0, logits, scratch, scratch, scratch + offsets[0], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(160u, status, descriptors, params + 1, intrinsic_1, scratch, scratch, scratch + offsets[1], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(144u, status, descriptors, params + 2, current_0, scratch, scratch, scratch + offsets[2], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(129u, status, descriptors, params + 3, scratch, scratch, scratch, scratch + offsets[3], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(32u, status, descriptors, params + 4, scratch + offsets[2], scratch + offsets[0], scratch + offsets[3], scratch + offsets[4], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(51u, status, descriptors, params + 5, scratch + offsets[4], scratch, scratch, scratch + offsets[5], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(144u, status, descriptors, params + 6, current_1, scratch, scratch, scratch + offsets[6], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(100u, status, descriptors, params + 7, scratch, scratch, scratch, scratch + offsets[7], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(96u, status, descriptors, params + 8, scratch + offsets[5], scratch + offsets[7], scratch, scratch + offsets[8], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(24u, status, descriptors, params + 9, scratch + offsets[6], scratch + offsets[8], scratch, scratch + offsets[9], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(129u, status, descriptors, params + 10, scratch, scratch, scratch, scratch + offsets[10], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(129u, status, descriptors, params + 11, scratch, scratch, scratch, scratch + offsets[11], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(32u, status, descriptors, params + 12, scratch + offsets[9], scratch + offsets[10], scratch + offsets[11], scratch + offsets[12], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== mtp_verify_tail#0 fused#1
function_name: ptir_m2_963137d4ffed0445_r1
region: library=1 library_op=3 schedule=3 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41704 prefix=@runtime tail_bytes=1205
kernel void ptir_m2_963137d4ffed0445_r1(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(65u, status, descriptors, params + 13, scratch + offsets[12], scratch, scratch, scratch + offsets[13], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== mtp_verify_tail#0 fused#2
function_name: ptir_m2_963137d4ffed0445_r2
region: library=0 library_op=0 schedule=1 nodes=11 inputs=3 outputs=0 sinks=3
ok: true
source: bytes=43519 prefix=@runtime tail_bytes=3020
kernel void ptir_m2_963137d4ffed0445_r2(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(48u, status, descriptors, params + 14, scratch + offsets[13], scratch, scratch, scratch + offsets[14], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(7u, status, descriptors, params + 15, scratch + offsets[14], scratch, scratch, scratch + offsets[15], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(56u, status, descriptors, params + 16, scratch + offsets[15], scratch, scratch, scratch + offsets[16], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(100u, status, descriptors, params + 17, scratch, scratch, scratch, scratch + offsets[17], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(23u, status, descriptors, params + 18, scratch + offsets[16], scratch + offsets[17], scratch, scratch + offsets[18], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(129u, status, descriptors, params + 19, scratch, scratch, scratch, scratch + offsets[19], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(32u, status, descriptors, params + 20, scratch + offsets[18], scratch + offsets[5], scratch + offsets[19], scratch + offsets[20], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(146u, status, descriptors, params + 21, scratch + offsets[20], scratch, scratch, pending_2, scratch, temporary);
  if (status->state != 1) return;
  current_2 = pending_2;
  ptir_m1_execute(51u, status, descriptors, params + 22, scratch + offsets[1], scratch, scratch, scratch + offsets[21], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(146u, status, descriptors, params + 23, scratch + offsets[21], scratch, scratch, pending_1, scratch, temporary);
  if (status->state != 1) return;
  current_1 = pending_1;
  ptir_m1_execute(146u, status, descriptors, params + 24, scratch + offsets[21], scratch, scratch, pending_3, scratch, temporary);
  if (status->state != 1) return;
  current_3 = pending_3;
}
=== end
=== nucleus_sample#0 singleton#0
function_name: ptir_m2_82921040f7bc2612_r0
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41545 prefix=@runtime tail_bytes=1046
kernel void ptir_m2_82921040f7bc2612_r0(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  ptir_m1_execute(160u, status, descriptors, params + 0, logits, scratch, scratch, scratch + offsets[0], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== nucleus_sample#0 singleton#1
function_name: ptir_m2_82921040f7bc2612_r1
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41548 prefix=@runtime tail_bytes=1049
kernel void ptir_m2_82921040f7bc2612_r1(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  ptir_m1_execute(144u, status, descriptors, params + 1, current_0, scratch, scratch, scratch + offsets[1], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== nucleus_sample#0 singleton#2
function_name: ptir_m2_82921040f7bc2612_r2
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41548 prefix=@runtime tail_bytes=1049
kernel void ptir_m2_82921040f7bc2612_r2(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  ptir_m1_execute(145u, status, descriptors, params + 2, current_1, scratch, scratch, scratch + offsets[2], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== nucleus_sample#0 singleton#3
function_name: ptir_m2_82921040f7bc2612_r3
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41558 prefix=@runtime tail_bytes=1059
kernel void ptir_m2_82921040f7bc2612_r3(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  ptir_m1_execute(49u, status, descriptors, params + 3, scratch + offsets[0], scratch, scratch, scratch + offsets[3], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== nucleus_sample#0 singleton#4
function_name: ptir_m2_82921040f7bc2612_r4
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41558 prefix=@runtime tail_bytes=1059
kernel void ptir_m2_82921040f7bc2612_r4(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  ptir_m1_execute(56u, status, descriptors, params + 4, scratch + offsets[3], scratch, scratch, scratch + offsets[4], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== nucleus_sample#0 singleton#5
function_name: ptir_m2_82921040f7bc2612_r5
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: true
source: bytes=41571 prefix=@runtime tail_bytes=1072
kernel void ptir_m2_82921040f7bc2612_r5(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  ptir_m1_execute(17u, status, descriptors, params + 5, scratch + offsets[0], scratch + offsets[4], scratch, scratch + offsets[5], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== nucleus_sample#0 singleton#6
function_name: ptir_m2_82921040f7bc2612_r6
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41557 prefix=@runtime tail_bytes=1058
kernel void ptir_m2_82921040f7bc2612_r6(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  ptir_m1_execute(1u, status, descriptors, params + 6, scratch + offsets[5], scratch, scratch, scratch + offsets[6], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== nucleus_sample#0 singleton#7
function_name: ptir_m2_82921040f7bc2612_r7
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41558 prefix=@runtime tail_bytes=1059
kernel void ptir_m2_82921040f7bc2612_r7(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  ptir_m1_execute(48u, status, descriptors, params + 7, scratch + offsets[6], scratch, scratch, scratch + offsets[7], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== nucleus_sample#0 singleton#9
function_name: ptir_m2_82921040f7bc2612_r9
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: true
source: bytes=41571 prefix=@runtime tail_bytes=1072
kernel void ptir_m2_82921040f7bc2612_r9(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  ptir_m1_execute(19u, status, descriptors, params + 9, scratch + offsets[6], scratch + offsets[8], scratch, scratch + offsets[9], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== nucleus_sample#0 singleton#10
function_name: ptir_m2_82921040f7bc2612_r10
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: true
source: bytes=41574 prefix=@runtime tail_bytes=1075
kernel void ptir_m2_82921040f7bc2612_r10(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  ptir_m1_execute(88u, status, descriptors, params + 10, scratch + offsets[9], scratch + offsets[2], scratch, scratch + offsets[10], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== nucleus_sample#0 singleton#11
function_name: ptir_m2_82921040f7bc2612_r11
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41549 prefix=@runtime tail_bytes=1050
kernel void ptir_m2_82921040f7bc2612_r11(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  ptir_m1_execute(129u, status, descriptors, params + 11, scratch, scratch, scratch, scratch + offsets[11], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== nucleus_sample#0 singleton#12
function_name: ptir_m2_82921040f7bc2612_r12
region: library=0 library_op=0 schedule=1 nodes=1 inputs=3 outputs=1 sinks=0
ok: true
source: bytes=41589 prefix=@runtime tail_bytes=1090
kernel void ptir_m2_82921040f7bc2612_r12(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  ptir_m1_execute(32u, status, descriptors, params + 12, scratch + offsets[10], scratch + offsets[0], scratch + offsets[11], scratch + offsets[12], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== nucleus_sample#0 singleton#13
function_name: ptir_m2_82921040f7bc2612_r13
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41562 prefix=@runtime tail_bytes=1063
kernel void ptir_m2_82921040f7bc2612_r13(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  ptir_m1_execute(113u, status, descriptors, params + 13, scratch + offsets[1], scratch, scratch, scratch + offsets[13], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== nucleus_sample#0 singleton#14
function_name: ptir_m2_82921040f7bc2612_r14
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: true
source: bytes=41576 prefix=@runtime tail_bytes=1077
kernel void ptir_m2_82921040f7bc2612_r14(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  ptir_m1_execute(16u, status, descriptors, params + 14, scratch + offsets[12], scratch + offsets[13], scratch, scratch + offsets[14], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== nucleus_sample#0 singleton#15
function_name: ptir_m2_82921040f7bc2612_r15
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41562 prefix=@runtime tail_bytes=1063
kernel void ptir_m2_82921040f7bc2612_r15(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  ptir_m1_execute(51u, status, descriptors, params + 15, scratch + offsets[14], scratch, scratch, scratch + offsets[15], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== nucleus_sample#0 singleton#16
function_name: ptir_m2_82921040f7bc2612_r16
region: library=0 library_op=0 schedule=0 nodes=1 inputs=1 outputs=0 sinks=1
ok: true
source: bytes=41576 prefix=@runtime tail_bytes=1077
kernel void ptir_m2_82921040f7bc2612_r16(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  ptir_m1_execute(146u, status, descriptors, params + 16, scratch + offsets[15], scratch, scratch, pending_2, scratch, temporary);
  if (status->state != 1) return;
  current_2 = pending_2;
}
=== end
=== nucleus_sample#0 fused#0
function_name: ptir_m2_82921040f7bc2612_r0
region: library=0 library_op=0 schedule=1 nodes=3 inputs=0 outputs=3 sinks=0
ok: true
source: bytes=41871 prefix=@runtime tail_bytes=1372
kernel void ptir_m2_82921040f7bc2612_r0(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  ptir_m1_execute(160u, status, descriptors, params + 0, logits, scratch, scratch, scratch + offsets[0], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(144u, status, descriptors, params + 1, current_0, scratch, scratch, scratch + offsets[1], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(145u, status, descriptors, params + 2, current_1, scratch, scratch, scratch + offsets[2], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== nucleus_sample#0 fused#1
function_name: ptir_m2_82921040f7bc2612_r1
region: library=1 library_op=0 schedule=3 nodes=13 inputs=3 outputs=1 sinks=0
ok: true
source: bytes=43717 prefix=@runtime tail_bytes=3218
kernel void ptir_m2_82921040f7bc2612_r1(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  ptir_m1_execute(49u, status, descriptors, params + 3, scratch + offsets[0], scratch, scratch, scratch + offsets[3], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(56u, status, descriptors, params + 4, scratch + offsets[3], scratch, scratch, scratch + offsets[4], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(17u, status, descriptors, params + 5, scratch + offsets[0], scratch + offsets[4], scratch, scratch + offsets[5], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(1u, status, descriptors, params + 6, scratch + offsets[5], scratch, scratch, scratch + offsets[6], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(48u, status, descriptors, params + 7, scratch + offsets[6], scratch, scratch, scratch + offsets[7], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(56u, status, descriptors, params + 8, scratch + offsets[7], scratch, scratch, scratch + offsets[8], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(19u, status, descriptors, params + 9, scratch + offsets[6], scratch + offsets[8], scratch, scratch + offsets[9], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(88u, status, descriptors, params + 10, scratch + offsets[9], scratch + offsets[2], scratch, scratch + offsets[10], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(129u, status, descriptors, params + 11, scratch, scratch, scratch, scratch + offsets[11], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(32u, status, descriptors, params + 12, scratch + offsets[10], scratch + offsets[0], scratch + offsets[11], scratch + offsets[12], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(113u, status, descriptors, params + 13, scratch + offsets[1], scratch, scratch, scratch + offsets[13], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(16u, status, descriptors, params + 14, scratch + offsets[12], scratch + offsets[13], scratch, scratch + offsets[14], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(51u, status, descriptors, params + 15, scratch + offsets[14], scratch, scratch, scratch + offsets[15], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== nucleus_sample#0 fused#2
function_name: ptir_m2_82921040f7bc2612_r2
region: library=0 library_op=0 schedule=0 nodes=1 inputs=1 outputs=0 sinks=1
ok: true
source: bytes=41575 prefix=@runtime tail_bytes=1076
kernel void ptir_m2_82921040f7bc2612_r2(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  ptir_m1_execute(146u, status, descriptors, params + 16, scratch + offsets[15], scratch, scratch, pending_2, scratch, temporary);
  if (status->state != 1) return;
  current_2 = pending_2;
}
=== end
=== pentathlon_iter#0 singleton#0
function_name: ptir_m2_14f667a41eb30669_r0
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: Metal binds only the logits buffer for intrinsics; intrinsic id 3 has no binding
=== end
=== pentathlon_iter#0 singleton#1
function_name: ptir_m2_14f667a41eb30669_r1
region: library=1 library_op=5 schedule=3 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41274 prefix=@runtime tail_bytes=775
kernel void ptir_m2_14f667a41eb30669_r1(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  ptir_m1_execute(161u, status, descriptors, params + 1, scratch + offsets[0], scratch, scratch, scratch + offsets[1], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== pentathlon_iter#0 singleton#2
function_name: ptir_m2_14f667a41eb30669_r2
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41263 prefix=@runtime tail_bytes=764
kernel void ptir_m2_14f667a41eb30669_r2(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  ptir_m1_execute(145u, status, descriptors, params + 2, current_0, scratch, scratch, scratch + offsets[2], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== pentathlon_iter#0 singleton#3
function_name: ptir_m2_14f667a41eb30669_r3
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41273 prefix=@runtime tail_bytes=774
kernel void ptir_m2_14f667a41eb30669_r3(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  ptir_m1_execute(57u, status, descriptors, params + 3, scratch + offsets[2], scratch, scratch, scratch + offsets[3], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== pentathlon_iter#0 singleton#4
function_name: ptir_m2_14f667a41eb30669_r4
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: true
source: bytes=41286 prefix=@runtime tail_bytes=787
kernel void ptir_m2_14f667a41eb30669_r4(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  ptir_m1_execute(88u, status, descriptors, params + 4, scratch + offsets[1], scratch + offsets[3], scratch, scratch + offsets[4], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== pentathlon_iter#0 singleton#5
function_name: ptir_m2_14f667a41eb30669_r5
region: library=1 library_op=5 schedule=3 nodes=1 inputs=1 outputs=0 sinks=0
ok: true
source: bytes=41261 prefix=@runtime tail_bytes=762
kernel void ptir_m2_14f667a41eb30669_r5(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  ptir_m1_execute(162u, status, descriptors, params + 5, scratch + offsets[4], scratch, scratch, scratch, scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== pentathlon_iter#0 fused#0
function_name: ptir_m2_14f667a41eb30669_r0
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: Metal binds only the logits buffer for intrinsics; intrinsic id 3 has no binding
=== end
=== pentathlon_iter#0 fused#1
function_name: ptir_m2_14f667a41eb30669_r1
region: library=1 library_op=5 schedule=3 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41274 prefix=@runtime tail_bytes=775
kernel void ptir_m2_14f667a41eb30669_r1(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  ptir_m1_execute(161u, status, descriptors, params + 1, scratch + offsets[0], scratch, scratch, scratch + offsets[1], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== pentathlon_iter#0 fused#2
function_name: ptir_m2_14f667a41eb30669_r2
region: library=0 library_op=0 schedule=1 nodes=3 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41622 prefix=@runtime tail_bytes=1123
kernel void ptir_m2_14f667a41eb30669_r2(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  ptir_m1_execute(145u, status, descriptors, params + 2, current_0, scratch, scratch, scratch + offsets[2], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(57u, status, descriptors, params + 3, scratch + offsets[2], scratch, scratch, scratch + offsets[3], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(88u, status, descriptors, params + 4, scratch + offsets[1], scratch + offsets[3], scratch, scratch + offsets[4], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== pentathlon_iter#0 fused#3
function_name: ptir_m2_14f667a41eb30669_r3
region: library=1 library_op=5 schedule=3 nodes=1 inputs=1 outputs=0 sinks=0
ok: true
source: bytes=41261 prefix=@runtime tail_bytes=762
kernel void ptir_m2_14f667a41eb30669_r3(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  ptir_m1_execute(162u, status, descriptors, params + 5, scratch + offsets[4], scratch, scratch, scratch, scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== pentathlon_iter#1 singleton#0
function_name: ptir_m2_dff3bc6795a9dcdf_r0
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41831 prefix=@runtime tail_bytes=1332
kernel void ptir_m2_dff3bc6795a9dcdf_r0(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  ptir_m1_execute(160u, status, descriptors, params + 0, logits, scratch, scratch, scratch + offsets[0], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== pentathlon_iter#1 singleton#1
function_name: ptir_m2_dff3bc6795a9dcdf_r1
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41834 prefix=@runtime tail_bytes=1335
kernel void ptir_m2_dff3bc6795a9dcdf_r1(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  ptir_m1_execute(144u, status, descriptors, params + 1, current_0, scratch, scratch, scratch + offsets[1], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== pentathlon_iter#1 singleton#3
function_name: ptir_m2_dff3bc6795a9dcdf_r3
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41844 prefix=@runtime tail_bytes=1345
kernel void ptir_m2_dff3bc6795a9dcdf_r3(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  ptir_m1_execute(49u, status, descriptors, params + 3, scratch + offsets[0], scratch, scratch, scratch + offsets[3], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== pentathlon_iter#1 singleton#4
function_name: ptir_m2_dff3bc6795a9dcdf_r4
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41844 prefix=@runtime tail_bytes=1345
kernel void ptir_m2_dff3bc6795a9dcdf_r4(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  ptir_m1_execute(56u, status, descriptors, params + 4, scratch + offsets[3], scratch, scratch, scratch + offsets[4], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== pentathlon_iter#1 singleton#5
function_name: ptir_m2_dff3bc6795a9dcdf_r5
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: true
source: bytes=41857 prefix=@runtime tail_bytes=1358
kernel void ptir_m2_dff3bc6795a9dcdf_r5(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  ptir_m1_execute(17u, status, descriptors, params + 5, scratch + offsets[0], scratch + offsets[4], scratch, scratch + offsets[5], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== pentathlon_iter#1 singleton#6
function_name: ptir_m2_dff3bc6795a9dcdf_r6
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41843 prefix=@runtime tail_bytes=1344
kernel void ptir_m2_dff3bc6795a9dcdf_r6(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  ptir_m1_execute(1u, status, descriptors, params + 6, scratch + offsets[5], scratch, scratch, scratch + offsets[6], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== pentathlon_iter#1 singleton#7
function_name: ptir_m2_dff3bc6795a9dcdf_r7
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41844 prefix=@runtime tail_bytes=1345
kernel void ptir_m2_dff3bc6795a9dcdf_r7(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  ptir_m1_execute(48u, status, descriptors, params + 7, scratch + offsets[6], scratch, scratch, scratch + offsets[7], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== pentathlon_iter#1 singleton#8
function_name: ptir_m2_dff3bc6795a9dcdf_r8
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41843 prefix=@runtime tail_bytes=1344
kernel void ptir_m2_dff3bc6795a9dcdf_r8(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  ptir_m1_execute(2u, status, descriptors, params + 8, scratch + offsets[7], scratch, scratch, scratch + offsets[8], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== pentathlon_iter#1 singleton#19
function_name: ptir_m2_dff3bc6795a9dcdf_r19
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41835 prefix=@runtime tail_bytes=1336
kernel void ptir_m2_dff3bc6795a9dcdf_r19(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  ptir_m1_execute(129u, status, descriptors, params + 19, scratch, scratch, scratch, scratch + offsets[19], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== pentathlon_iter#1 singleton#20
function_name: ptir_m2_dff3bc6795a9dcdf_r20
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: true
source: bytes=41862 prefix=@runtime tail_bytes=1363
kernel void ptir_m2_dff3bc6795a9dcdf_r20(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  ptir_m1_execute(18u, status, descriptors, params + 20, scratch + offsets[18], scratch + offsets[19], scratch, scratch + offsets[20], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== pentathlon_iter#1 singleton#23
function_name: ptir_m2_dff3bc6795a9dcdf_r23
region: library=0 library_op=0 schedule=1 nodes=1 inputs=3 outputs=1 sinks=0
ok: true
source: bytes=41875 prefix=@runtime tail_bytes=1376
kernel void ptir_m2_dff3bc6795a9dcdf_r23(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  ptir_m1_execute(32u, status, descriptors, params + 23, scratch + offsets[2], scratch + offsets[10], scratch + offsets[22], scratch + offsets[23], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== pentathlon_iter#1 singleton#25
function_name: ptir_m2_dff3bc6795a9dcdf_r25
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41848 prefix=@runtime tail_bytes=1349
kernel void ptir_m2_dff3bc6795a9dcdf_r25(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  ptir_m1_execute(57u, status, descriptors, params + 25, scratch + offsets[24], scratch, scratch, scratch + offsets[25], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== pentathlon_iter#1 singleton#28
function_name: ptir_m2_dff3bc6795a9dcdf_r28
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: true
source: bytes=41862 prefix=@runtime tail_bytes=1363
kernel void ptir_m2_dff3bc6795a9dcdf_r28(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  ptir_m1_execute(16u, status, descriptors, params + 28, scratch + offsets[26], scratch + offsets[27], scratch, scratch + offsets[28], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== pentathlon_iter#1 singleton#29
function_name: ptir_m2_dff3bc6795a9dcdf_r29
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: true
source: bytes=41862 prefix=@runtime tail_bytes=1363
kernel void ptir_m2_dff3bc6795a9dcdf_r29(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  ptir_m1_execute(23u, status, descriptors, params + 29, scratch + offsets[23], scratch + offsets[28], scratch, scratch + offsets[29], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== pentathlon_iter#1 singleton#31
function_name: ptir_m2_dff3bc6795a9dcdf_r31
region: library=1 library_op=1 schedule=3 nodes=1 inputs=1 outputs=2 sinks=0
ok: true
source: bytes=41862 prefix=@runtime tail_bytes=1363
kernel void ptir_m2_dff3bc6795a9dcdf_r31(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  ptir_m1_execute(81u, status, descriptors, params + 31, scratch + offsets[30], scratch, scratch, scratch + offsets[31], scratch + offsets[32], temporary);
  if (status->state != 1) return;
}
=== end
=== pentathlon_iter#1 singleton#34
function_name: ptir_m2_dff3bc6795a9dcdf_r34
region: library=0 library_op=0 schedule=0 nodes=1 inputs=1 outputs=0 sinks=1
ok: true
source: bytes=41862 prefix=@runtime tail_bytes=1363
kernel void ptir_m2_dff3bc6795a9dcdf_r34(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  ptir_m1_execute(146u, status, descriptors, params + 34, scratch + offsets[33], scratch, scratch, pending_2, scratch, temporary);
  if (status->state != 1) return;
  current_2 = pending_2;
}
=== end
=== pentathlon_iter#1 fused#0
function_name: ptir_m2_dff3bc6795a9dcdf_r0
region: library=0 library_op=0 schedule=1 nodes=31 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=47183 prefix=@runtime tail_bytes=6684
kernel void ptir_m2_dff3bc6795a9dcdf_r0(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  ptir_m1_execute(160u, status, descriptors, params + 0, logits, scratch, scratch, scratch + offsets[0], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(144u, status, descriptors, params + 1, current_0, scratch, scratch, scratch + offsets[1], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(144u, status, descriptors, params + 2, current_1, scratch, scratch, scratch + offsets[2], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(49u, status, descriptors, params + 3, scratch + offsets[0], scratch, scratch, scratch + offsets[3], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(56u, status, descriptors, params + 4, scratch + offsets[3], scratch, scratch, scratch + offsets[4], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(17u, status, descriptors, params + 5, scratch + offsets[0], scratch + offsets[4], scratch, scratch + offsets[5], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(1u, status, descriptors, params + 6, scratch + offsets[5], scratch, scratch, scratch + offsets[6], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(48u, status, descriptors, params + 7, scratch + offsets[6], scratch, scratch, scratch + offsets[7], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(2u, status, descriptors, params + 8, scratch + offsets[7], scratch, scratch, scratch + offsets[8], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(56u, status, descriptors, params + 9, scratch + offsets[8], scratch, scratch, scratch + offsets[9], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(17u, status, descriptors, params + 10, scratch + offsets[5], scratch + offsets[9], scratch, scratch + offsets[10], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(49u, status, descriptors, params + 11, scratch + offsets[1], scratch, scratch, scratch + offsets[11], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(56u, status, descriptors, params + 12, scratch + offsets[11], scratch, scratch, scratch + offsets[12], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(17u, status, descriptors, params + 13, scratch + offsets[1], scratch + offsets[12], scratch, scratch + offsets[13], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(1u, status, descriptors, params + 14, scratch + offsets[13], scratch, scratch, scratch + offsets[14], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(48u, status, descriptors, params + 15, scratch + offsets[14], scratch, scratch, scratch + offsets[15], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(2u, status, descriptors, params + 16, scratch + offsets[15], scratch, scratch, scratch + offsets[16], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(56u, status, descriptors, params + 17, scratch + offsets[16], scratch, scratch, scratch + offsets[17], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(17u, status, descriptors, params + 18, scratch + offsets[13], scratch + offsets[17], scratch, scratch + offsets[18], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(129u, status, descriptors, params + 19, scratch, scratch, scratch, scratch + offsets[19], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(18u, status, descriptors, params + 20, scratch + offsets[18], scratch + offsets[19], scratch, scratch + offsets[20], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(17u, status, descriptors, params + 21, scratch + offsets[10], scratch + offsets[20], scratch, scratch + offsets[21], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(129u, status, descriptors, params + 22, scratch, scratch, scratch, scratch + offsets[22], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(32u, status, descriptors, params + 23, scratch + offsets[2], scratch + offsets[10], scratch + offsets[22], scratch + offsets[23], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(49u, status, descriptors, params + 24, scratch + offsets[23], scratch, scratch, scratch + offsets[24], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(57u, status, descriptors, params + 25, scratch + offsets[24], scratch, scratch, scratch + offsets[25], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(56u, status, descriptors, params + 26, scratch + offsets[25], scratch, scratch, scratch + offsets[26], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(129u, status, descriptors, params + 27, scratch, scratch, scratch, scratch + offsets[27], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(16u, status, descriptors, params + 28, scratch + offsets[26], scratch + offsets[27], scratch, scratch + offsets[28], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(23u, status, descriptors, params + 29, scratch + offsets[23], scratch + offsets[28], scratch, scratch + offsets[29], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(32u, status, descriptors, params + 30, scratch + offsets[29], scratch + offsets[21], scratch + offsets[22], scratch + offsets[30], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== pentathlon_iter#1 fused#1
function_name: ptir_m2_dff3bc6795a9dcdf_r1
region: library=1 library_op=1 schedule=3 nodes=1 inputs=1 outputs=2 sinks=0
ok: true
source: bytes=41861 prefix=@runtime tail_bytes=1362
kernel void ptir_m2_dff3bc6795a9dcdf_r1(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  ptir_m1_execute(81u, status, descriptors, params + 31, scratch + offsets[30], scratch, scratch, scratch + offsets[31], scratch + offsets[32], temporary);
  if (status->state != 1) return;
}
=== end
=== pentathlon_iter#1 fused#2
function_name: ptir_m2_dff3bc6795a9dcdf_r2
region: library=0 library_op=0 schedule=1 nodes=6 inputs=2 outputs=0 sinks=3
ok: false
error: Metal binds only the logits buffer for intrinsics; intrinsic id 4 has no binding
=== end
=== pivot_predicates_multistage#1 singleton#0
function_name: ptir_m2_09a2719886970a85_r0
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41688 prefix=@runtime tail_bytes=1189
kernel void ptir_m2_09a2719886970a85_r0(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(160u, status, descriptors, params + 0, logits, scratch, scratch, scratch + offsets[0], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== pivot_predicates_multistage#1 singleton#1
function_name: ptir_m2_09a2719886970a85_r1
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41701 prefix=@runtime tail_bytes=1202
kernel void ptir_m2_09a2719886970a85_r1(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(49u, status, descriptors, params + 1, scratch + offsets[0], scratch, scratch, scratch + offsets[1], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== pivot_predicates_multistage#1 singleton#2
function_name: ptir_m2_09a2719886970a85_r2
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41701 prefix=@runtime tail_bytes=1202
kernel void ptir_m2_09a2719886970a85_r2(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(56u, status, descriptors, params + 2, scratch + offsets[1], scratch, scratch, scratch + offsets[2], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== pivot_predicates_multistage#1 singleton#3
function_name: ptir_m2_09a2719886970a85_r3
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: true
source: bytes=41714 prefix=@runtime tail_bytes=1215
kernel void ptir_m2_09a2719886970a85_r3(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(17u, status, descriptors, params + 3, scratch + offsets[0], scratch + offsets[2], scratch, scratch + offsets[3], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== pivot_predicates_multistage#1 singleton#4
function_name: ptir_m2_09a2719886970a85_r4
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41700 prefix=@runtime tail_bytes=1201
kernel void ptir_m2_09a2719886970a85_r4(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(1u, status, descriptors, params + 4, scratch + offsets[3], scratch, scratch, scratch + offsets[4], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== pivot_predicates_multistage#1 singleton#5
function_name: ptir_m2_09a2719886970a85_r5
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41701 prefix=@runtime tail_bytes=1202
kernel void ptir_m2_09a2719886970a85_r5(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(48u, status, descriptors, params + 5, scratch + offsets[4], scratch, scratch, scratch + offsets[5], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== pivot_predicates_multistage#1 singleton#7
function_name: ptir_m2_09a2719886970a85_r7
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: true
source: bytes=41714 prefix=@runtime tail_bytes=1215
kernel void ptir_m2_09a2719886970a85_r7(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(19u, status, descriptors, params + 7, scratch + offsets[4], scratch + offsets[6], scratch, scratch + offsets[7], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== pivot_predicates_multistage#1 singleton#8
function_name: ptir_m2_09a2719886970a85_r8
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41691 prefix=@runtime tail_bytes=1192
kernel void ptir_m2_09a2719886970a85_r8(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(145u, status, descriptors, params + 8, current_0, scratch, scratch, scratch + offsets[8], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== pivot_predicates_multistage#1 singleton#10
function_name: ptir_m2_09a2719886970a85_r10
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: true
source: bytes=41717 prefix=@runtime tail_bytes=1218
kernel void ptir_m2_09a2719886970a85_r10(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(88u, status, descriptors, params + 10, scratch + offsets[7], scratch + offsets[8], scratch, scratch + offsets[10], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== pivot_predicates_multistage#1 singleton#12
function_name: ptir_m2_09a2719886970a85_r12
region: library=0 library_op=0 schedule=0 nodes=1 inputs=1 outputs=0 sinks=1
ok: true
source: bytes=41719 prefix=@runtime tail_bytes=1220
kernel void ptir_m2_09a2719886970a85_r12(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(146u, status, descriptors, params + 12, scratch + offsets[10], scratch, scratch, pending_2, scratch, temporary);
  if (status->state != 1) return;
  current_2 = pending_2;
}
=== end
=== pivot_predicates_multistage#1 fused#0
function_name: ptir_m2_09a2719886970a85_r0
region: library=0 library_op=0 schedule=1 nodes=14 inputs=0 outputs=0 sinks=2
ok: true
source: bytes=44006 prefix=@runtime tail_bytes=3507
kernel void ptir_m2_09a2719886970a85_r0(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  ptir_m1_execute(160u, status, descriptors, params + 0, logits, scratch, scratch, scratch + offsets[0], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(49u, status, descriptors, params + 1, scratch + offsets[0], scratch, scratch, scratch + offsets[1], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(56u, status, descriptors, params + 2, scratch + offsets[1], scratch, scratch, scratch + offsets[2], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(17u, status, descriptors, params + 3, scratch + offsets[0], scratch + offsets[2], scratch, scratch + offsets[3], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(1u, status, descriptors, params + 4, scratch + offsets[3], scratch, scratch, scratch + offsets[4], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(48u, status, descriptors, params + 5, scratch + offsets[4], scratch, scratch, scratch + offsets[5], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(56u, status, descriptors, params + 6, scratch + offsets[5], scratch, scratch, scratch + offsets[6], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(19u, status, descriptors, params + 7, scratch + offsets[4], scratch + offsets[6], scratch, scratch + offsets[7], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(145u, status, descriptors, params + 8, current_0, scratch, scratch, scratch + offsets[8], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(145u, status, descriptors, params + 9, current_1, scratch, scratch, scratch + offsets[9], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(88u, status, descriptors, params + 10, scratch + offsets[7], scratch + offsets[8], scratch, scratch + offsets[10], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(88u, status, descriptors, params + 11, scratch + offsets[7], scratch + offsets[9], scratch, scratch + offsets[11], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(146u, status, descriptors, params + 12, scratch + offsets[10], scratch, scratch, pending_2, scratch, temporary);
  if (status->state != 1) return;
  current_2 = pending_2;
  ptir_m1_execute(146u, status, descriptors, params + 13, scratch + offsets[11], scratch, scratch, pending_3, scratch, temporary);
  if (status->state != 1) return;
  current_3 = pending_3;
}
=== end
=== section3_masked_gumbel#0 singleton#0
function_name: ptir_m2_4cf3882bf2fff913_r0
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41831 prefix=@runtime tail_bytes=1332
kernel void ptir_m2_4cf3882bf2fff913_r0(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  ptir_m1_execute(160u, status, descriptors, params + 0, logits, scratch, scratch, scratch + offsets[0], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== section3_masked_gumbel#0 singleton#1
function_name: ptir_m2_4cf3882bf2fff913_r1
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41844 prefix=@runtime tail_bytes=1345
kernel void ptir_m2_4cf3882bf2fff913_r1(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  ptir_m1_execute(57u, status, descriptors, params + 1, scratch + offsets[0], scratch, scratch, scratch + offsets[1], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== section3_masked_gumbel#0 singleton#2
function_name: ptir_m2_4cf3882bf2fff913_r2
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41834 prefix=@runtime tail_bytes=1335
kernel void ptir_m2_4cf3882bf2fff913_r2(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  ptir_m1_execute(144u, status, descriptors, params + 2, current_0, scratch, scratch, scratch + offsets[2], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== section3_masked_gumbel#0 singleton#4
function_name: ptir_m2_4cf3882bf2fff913_r4
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41845 prefix=@runtime tail_bytes=1346
kernel void ptir_m2_4cf3882bf2fff913_r4(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  ptir_m1_execute(113u, status, descriptors, params + 4, scratch + offsets[2], scratch, scratch, scratch + offsets[4], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== section3_masked_gumbel#0 singleton#5
function_name: ptir_m2_4cf3882bf2fff913_r5
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41832 prefix=@runtime tail_bytes=1333
kernel void ptir_m2_4cf3882bf2fff913_r5(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  ptir_m1_execute(129u, status, descriptors, params + 5, scratch, scratch, scratch, scratch + offsets[5], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== section3_masked_gumbel#0 singleton#6
function_name: ptir_m2_4cf3882bf2fff913_r6
region: library=0 library_op=0 schedule=1 nodes=1 inputs=3 outputs=1 sinks=0
ok: true
source: bytes=41870 prefix=@runtime tail_bytes=1371
kernel void ptir_m2_4cf3882bf2fff913_r6(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  ptir_m1_execute(32u, status, descriptors, params + 6, scratch + offsets[3], scratch + offsets[1], scratch + offsets[5], scratch + offsets[6], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== section3_masked_gumbel#0 singleton#7
function_name: ptir_m2_4cf3882bf2fff913_r7
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: true
source: bytes=41857 prefix=@runtime tail_bytes=1358
kernel void ptir_m2_4cf3882bf2fff913_r7(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  ptir_m1_execute(16u, status, descriptors, params + 7, scratch + offsets[6], scratch + offsets[4], scratch, scratch + offsets[7], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== section3_masked_gumbel#0 singleton#8
function_name: ptir_m2_4cf3882bf2fff913_r8
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41844 prefix=@runtime tail_bytes=1345
kernel void ptir_m2_4cf3882bf2fff913_r8(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  ptir_m1_execute(51u, status, descriptors, params + 8, scratch + offsets[7], scratch, scratch, scratch + offsets[8], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== section3_masked_gumbel#0 singleton#9
function_name: ptir_m2_4cf3882bf2fff913_r9
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41832 prefix=@runtime tail_bytes=1333
kernel void ptir_m2_4cf3882bf2fff913_r9(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  ptir_m1_execute(100u, status, descriptors, params + 9, scratch, scratch, scratch, scratch + offsets[9], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== section3_masked_gumbel#0 singleton#11
function_name: ptir_m2_4cf3882bf2fff913_r11
region: library=0 library_op=0 schedule=0 nodes=1 inputs=1 outputs=0 sinks=1
ok: true
source: bytes=41862 prefix=@runtime tail_bytes=1363
kernel void ptir_m2_4cf3882bf2fff913_r11(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  ptir_m1_execute(146u, status, descriptors, params + 11, scratch + offsets[10], scratch, scratch, pending_0, scratch, temporary);
  if (status->state != 1) return;
  current_0 = pending_0;
}
=== end
=== section3_masked_gumbel#0 fused#0
function_name: ptir_m2_4cf3882bf2fff913_r0
region: library=0 library_op=0 schedule=1 nodes=19 inputs=0 outputs=0 sinks=4
ok: true
source: bytes=45025 prefix=@runtime tail_bytes=4526
kernel void ptir_m2_4cf3882bf2fff913_r0(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  ptir_m1_execute(160u, status, descriptors, params + 0, logits, scratch, scratch, scratch + offsets[0], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(57u, status, descriptors, params + 1, scratch + offsets[0], scratch, scratch, scratch + offsets[1], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(144u, status, descriptors, params + 2, current_0, scratch, scratch, scratch + offsets[2], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(144u, status, descriptors, params + 3, current_1, scratch, scratch, scratch + offsets[3], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(113u, status, descriptors, params + 4, scratch + offsets[2], scratch, scratch, scratch + offsets[4], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(129u, status, descriptors, params + 5, scratch, scratch, scratch, scratch + offsets[5], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(32u, status, descriptors, params + 6, scratch + offsets[3], scratch + offsets[1], scratch + offsets[5], scratch + offsets[6], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(16u, status, descriptors, params + 7, scratch + offsets[6], scratch + offsets[4], scratch, scratch + offsets[7], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(51u, status, descriptors, params + 8, scratch + offsets[7], scratch, scratch, scratch + offsets[8], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(100u, status, descriptors, params + 9, scratch, scratch, scratch, scratch + offsets[9], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(16u, status, descriptors, params + 10, scratch + offsets[2], scratch + offsets[9], scratch, scratch + offsets[10], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(146u, status, descriptors, params + 11, scratch + offsets[10], scratch, scratch, pending_0, scratch, temporary);
  if (status->state != 1) return;
  current_0 = pending_0;
  ptir_m1_execute(57u, status, descriptors, params + 12, scratch + offsets[8], scratch, scratch, scratch + offsets[11], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(146u, status, descriptors, params + 13, scratch + offsets[11], scratch, scratch, pending_2, scratch, temporary);
  if (status->state != 1) return;
  current_2 = pending_2;
  ptir_m1_execute(144u, status, descriptors, params + 14, current_3, scratch, scratch, scratch + offsets[12], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(129u, status, descriptors, params + 15, scratch, scratch, scratch, scratch + offsets[13], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(16u, status, descriptors, params + 16, scratch + offsets[12], scratch + offsets[13], scratch, scratch + offsets[14], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(146u, status, descriptors, params + 17, scratch + offsets[14], scratch, scratch, pending_3, scratch, temporary);
  if (status->state != 1) return;
  current_3 = pending_3;
  ptir_m1_execute(146u, status, descriptors, params + 18, scratch + offsets[11], scratch, scratch, pending_4, scratch, temporary);
  if (status->state != 1) return;
  current_4 = pending_4;
}
=== end
=== structured_masks#0 singleton#0
function_name: ptir_m2_1752e0dfb0bf5827_r0
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=42120 prefix=@runtime tail_bytes=1621
kernel void ptir_m2_1752e0dfb0bf5827_r0(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    const device uchar* committed_5 [[buffer(17)]],
    device uchar* pending_5 [[buffer(18)]],
    const device uchar* committed_6 [[buffer(19)]],
    device uchar* pending_6 [[buffer(20)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  const device uchar* current_5 = committed_5;
  const device uchar* current_6 = committed_6;
  ptir_m1_execute(144u, status, descriptors, params + 0, current_0, scratch, scratch, scratch + offsets[0], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== structured_masks#0 singleton#3
function_name: ptir_m2_1752e0dfb0bf5827_r3
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=42131 prefix=@runtime tail_bytes=1632
kernel void ptir_m2_1752e0dfb0bf5827_r3(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    const device uchar* committed_5 [[buffer(17)]],
    device uchar* pending_5 [[buffer(18)]],
    const device uchar* committed_6 [[buffer(19)]],
    device uchar* pending_6 [[buffer(20)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  const device uchar* current_5 = committed_5;
  const device uchar* current_6 = committed_6;
  ptir_m1_execute(102u, status, descriptors, params + 3, scratch + offsets[0], scratch, scratch, scratch + offsets[3], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== structured_masks#0 singleton#4
function_name: ptir_m2_1752e0dfb0bf5827_r4
region: library=0 library_op=0 schedule=0 nodes=1 inputs=1 outputs=0 sinks=1
ok: true
source: bytes=42145 prefix=@runtime tail_bytes=1646
kernel void ptir_m2_1752e0dfb0bf5827_r4(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    const device uchar* committed_5 [[buffer(17)]],
    device uchar* pending_5 [[buffer(18)]],
    const device uchar* committed_6 [[buffer(19)]],
    device uchar* pending_6 [[buffer(20)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  const device uchar* current_5 = committed_5;
  const device uchar* current_6 = committed_6;
  ptir_m1_execute(146u, status, descriptors, params + 4, scratch + offsets[3], scratch, scratch, pending_3, scratch, temporary);
  if (status->state != 1) return;
  current_3 = pending_3;
}
=== end
=== structured_masks#0 singleton#5
function_name: ptir_m2_1752e0dfb0bf5827_r5
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=42131 prefix=@runtime tail_bytes=1632
kernel void ptir_m2_1752e0dfb0bf5827_r5(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    const device uchar* committed_5 [[buffer(17)]],
    device uchar* pending_5 [[buffer(18)]],
    const device uchar* committed_6 [[buffer(19)]],
    device uchar* pending_6 [[buffer(20)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  const device uchar* current_5 = committed_5;
  const device uchar* current_6 = committed_6;
  ptir_m1_execute(103u, status, descriptors, params + 5, scratch + offsets[0], scratch, scratch, scratch + offsets[4], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== structured_masks#0 singleton#7
function_name: ptir_m2_1752e0dfb0bf5827_r7
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=42131 prefix=@runtime tail_bytes=1632
kernel void ptir_m2_1752e0dfb0bf5827_r7(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    const device uchar* committed_5 [[buffer(17)]],
    device uchar* pending_5 [[buffer(18)]],
    const device uchar* committed_6 [[buffer(19)]],
    device uchar* pending_6 [[buffer(20)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  const device uchar* current_5 = committed_5;
  const device uchar* current_6 = committed_6;
  ptir_m1_execute(104u, status, descriptors, params + 7, scratch + offsets[0], scratch, scratch, scratch + offsets[5], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== structured_masks#0 singleton#9
function_name: ptir_m2_1752e0dfb0bf5827_r9
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=42130 prefix=@runtime tail_bytes=1631
kernel void ptir_m2_1752e0dfb0bf5827_r9(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    const device uchar* committed_5 [[buffer(17)]],
    device uchar* pending_5 [[buffer(18)]],
    const device uchar* committed_6 [[buffer(19)]],
    device uchar* pending_6 [[buffer(20)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  const device uchar* current_5 = committed_5;
  const device uchar* current_6 = committed_6;
  ptir_m1_execute(57u, status, descriptors, params + 9, scratch + offsets[1], scratch, scratch, scratch + offsets[6], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== structured_masks#0 singleton#10
function_name: ptir_m2_1752e0dfb0bf5827_r10
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=42120 prefix=@runtime tail_bytes=1621
kernel void ptir_m2_1752e0dfb0bf5827_r10(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    const device uchar* committed_5 [[buffer(17)]],
    device uchar* pending_5 [[buffer(18)]],
    const device uchar* committed_6 [[buffer(19)]],
    device uchar* pending_6 [[buffer(20)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  const device uchar* current_5 = committed_5;
  const device uchar* current_6 = committed_6;
  ptir_m1_execute(100u, status, descriptors, params + 10, scratch, scratch, scratch, scratch + offsets[7], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== structured_masks#0 singleton#11
function_name: ptir_m2_1752e0dfb0bf5827_r11
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=42120 prefix=@runtime tail_bytes=1621
kernel void ptir_m2_1752e0dfb0bf5827_r11(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    const device uchar* committed_5 [[buffer(17)]],
    device uchar* pending_5 [[buffer(18)]],
    const device uchar* committed_6 [[buffer(19)]],
    device uchar* pending_6 [[buffer(20)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  const device uchar* current_5 = committed_5;
  const device uchar* current_6 = committed_6;
  ptir_m1_execute(129u, status, descriptors, params + 11, scratch, scratch, scratch, scratch + offsets[8], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== structured_masks#0 singleton#12
function_name: ptir_m2_1752e0dfb0bf5827_r12
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: true
source: bytes=42145 prefix=@runtime tail_bytes=1646
kernel void ptir_m2_1752e0dfb0bf5827_r12(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    const device uchar* committed_5 [[buffer(17)]],
    device uchar* pending_5 [[buffer(18)]],
    const device uchar* committed_6 [[buffer(19)]],
    device uchar* pending_6 [[buffer(20)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  const device uchar* current_5 = committed_5;
  const device uchar* current_6 = committed_6;
  ptir_m1_execute(19u, status, descriptors, params + 12, scratch + offsets[7], scratch + offsets[8], scratch, scratch + offsets[9], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== structured_masks#0 singleton#14
function_name: ptir_m2_1752e0dfb0bf5827_r14
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: true
source: bytes=42147 prefix=@runtime tail_bytes=1648
kernel void ptir_m2_1752e0dfb0bf5827_r14(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    const device uchar* committed_5 [[buffer(17)]],
    device uchar* pending_5 [[buffer(18)]],
    const device uchar* committed_6 [[buffer(19)]],
    device uchar* pending_6 [[buffer(20)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  const device uchar* current_5 = committed_5;
  const device uchar* current_6 = committed_6;
  ptir_m1_execute(31u, status, descriptors, params + 14, scratch + offsets[7], scratch + offsets[10], scratch, scratch + offsets[11], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== structured_masks#0 singleton#15
function_name: ptir_m2_1752e0dfb0bf5827_r15
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: true
source: bytes=42147 prefix=@runtime tail_bytes=1648
kernel void ptir_m2_1752e0dfb0bf5827_r15(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    const device uchar* committed_5 [[buffer(17)]],
    device uchar* pending_5 [[buffer(18)]],
    const device uchar* committed_6 [[buffer(19)]],
    device uchar* pending_6 [[buffer(20)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  const device uchar* current_5 = committed_5;
  const device uchar* current_6 = committed_6;
  ptir_m1_execute(18u, status, descriptors, params + 15, scratch + offsets[9], scratch + offsets[10], scratch, scratch + offsets[12], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== structured_masks#0 singleton#16
function_name: ptir_m2_1752e0dfb0bf5827_r16
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: true
source: bytes=42148 prefix=@runtime tail_bytes=1649
kernel void ptir_m2_1752e0dfb0bf5827_r16(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    const device uchar* committed_5 [[buffer(17)]],
    device uchar* pending_5 [[buffer(18)]],
    const device uchar* committed_6 [[buffer(19)]],
    device uchar* pending_6 [[buffer(20)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  const device uchar* current_5 = committed_5;
  const device uchar* current_6 = committed_6;
  ptir_m1_execute(16u, status, descriptors, params + 16, scratch + offsets[11], scratch + offsets[12], scratch, scratch + offsets[13], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== structured_masks#0 singleton#17
function_name: ptir_m2_1752e0dfb0bf5827_r17
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: true
source: bytes=42147 prefix=@runtime tail_bytes=1648
kernel void ptir_m2_1752e0dfb0bf5827_r17(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    const device uchar* committed_5 [[buffer(17)]],
    device uchar* pending_5 [[buffer(18)]],
    const device uchar* committed_6 [[buffer(19)]],
    device uchar* pending_6 [[buffer(20)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  const device uchar* current_5 = committed_5;
  const device uchar* current_6 = committed_6;
  ptir_m1_execute(96u, status, descriptors, params + 17, scratch + offsets[6], scratch + offsets[13], scratch, scratch + offsets[14], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== structured_masks#0 singleton#24
function_name: ptir_m2_1752e0dfb0bf5827_r24
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: true
source: bytes=42148 prefix=@runtime tail_bytes=1649
kernel void ptir_m2_1752e0dfb0bf5827_r24(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    const device uchar* committed_5 [[buffer(17)]],
    device uchar* pending_5 [[buffer(18)]],
    const device uchar* committed_6 [[buffer(19)]],
    device uchar* pending_6 [[buffer(20)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  const device uchar* current_5 = committed_5;
  const device uchar* current_6 = committed_6;
  ptir_m1_execute(24u, status, descriptors, params + 24, scratch + offsets[19], scratch + offsets[20], scratch, scratch + offsets[21], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== structured_masks#0 singleton#25
function_name: ptir_m2_1752e0dfb0bf5827_r25
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=42133 prefix=@runtime tail_bytes=1634
kernel void ptir_m2_1752e0dfb0bf5827_r25(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    const device uchar* committed_5 [[buffer(17)]],
    device uchar* pending_5 [[buffer(18)]],
    const device uchar* committed_6 [[buffer(19)]],
    device uchar* pending_6 [[buffer(20)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  const device uchar* current_5 = committed_5;
  const device uchar* current_6 = committed_6;
  ptir_m1_execute(7u, status, descriptors, params + 25, scratch + offsets[21], scratch, scratch, scratch + offsets[22], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== structured_masks#0 singleton#26
function_name: ptir_m2_1752e0dfb0bf5827_r26
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=42134 prefix=@runtime tail_bytes=1635
kernel void ptir_m2_1752e0dfb0bf5827_r26(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    const device uchar* committed_5 [[buffer(17)]],
    device uchar* pending_5 [[buffer(18)]],
    const device uchar* committed_6 [[buffer(19)]],
    device uchar* pending_6 [[buffer(20)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  const device uchar* current_5 = committed_5;
  const device uchar* current_6 = committed_6;
  ptir_m1_execute(49u, status, descriptors, params + 26, scratch + offsets[22], scratch, scratch, scratch + offsets[23], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== structured_masks#0 fused#0
function_name: ptir_m2_1752e0dfb0bf5827_r0
region: library=0 library_op=0 schedule=1 nodes=29 inputs=0 outputs=0 sinks=4
ok: true
source: bytes=47125 prefix=@runtime tail_bytes=6626
kernel void ptir_m2_1752e0dfb0bf5827_r0(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    const device uchar* committed_3 [[buffer(13)]],
    device uchar* pending_3 [[buffer(14)]],
    const device uchar* committed_4 [[buffer(15)]],
    device uchar* pending_4 [[buffer(16)]],
    const device uchar* committed_5 [[buffer(17)]],
    device uchar* pending_5 [[buffer(18)]],
    const device uchar* committed_6 [[buffer(19)]],
    device uchar* pending_6 [[buffer(20)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  const device uchar* current_3 = committed_3;
  const device uchar* current_4 = committed_4;
  const device uchar* current_5 = committed_5;
  const device uchar* current_6 = committed_6;
  ptir_m1_execute(144u, status, descriptors, params + 0, current_0, scratch, scratch, scratch + offsets[0], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(144u, status, descriptors, params + 1, current_1, scratch, scratch, scratch + offsets[1], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(144u, status, descriptors, params + 2, current_2, scratch, scratch, scratch + offsets[2], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(102u, status, descriptors, params + 3, scratch + offsets[0], scratch, scratch, scratch + offsets[3], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(146u, status, descriptors, params + 4, scratch + offsets[3], scratch, scratch, pending_3, scratch, temporary);
  if (status->state != 1) return;
  current_3 = pending_3;
  ptir_m1_execute(103u, status, descriptors, params + 5, scratch + offsets[0], scratch, scratch, scratch + offsets[4], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(146u, status, descriptors, params + 6, scratch + offsets[4], scratch, scratch, pending_4, scratch, temporary);
  if (status->state != 1) return;
  current_4 = pending_4;
  ptir_m1_execute(104u, status, descriptors, params + 7, scratch + offsets[0], scratch, scratch, scratch + offsets[5], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(146u, status, descriptors, params + 8, scratch + offsets[5], scratch, scratch, pending_5, scratch, temporary);
  if (status->state != 1) return;
  current_5 = pending_5;
  ptir_m1_execute(57u, status, descriptors, params + 9, scratch + offsets[1], scratch, scratch, scratch + offsets[6], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(100u, status, descriptors, params + 10, scratch, scratch, scratch, scratch + offsets[7], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(129u, status, descriptors, params + 11, scratch, scratch, scratch, scratch + offsets[8], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(19u, status, descriptors, params + 12, scratch + offsets[7], scratch + offsets[8], scratch, scratch + offsets[9], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(129u, status, descriptors, params + 13, scratch, scratch, scratch, scratch + offsets[10], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(31u, status, descriptors, params + 14, scratch + offsets[7], scratch + offsets[10], scratch, scratch + offsets[11], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(18u, status, descriptors, params + 15, scratch + offsets[9], scratch + offsets[10], scratch, scratch + offsets[12], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(16u, status, descriptors, params + 16, scratch + offsets[11], scratch + offsets[12], scratch, scratch + offsets[13], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(96u, status, descriptors, params + 17, scratch + offsets[6], scratch + offsets[13], scratch, scratch + offsets[14], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(19u, status, descriptors, params + 18, scratch + offsets[7], scratch + offsets[10], scratch, scratch + offsets[15], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(129u, status, descriptors, params + 19, scratch, scratch, scratch, scratch + offsets[16], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(31u, status, descriptors, params + 20, scratch + offsets[15], scratch + offsets[16], scratch, scratch + offsets[17], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(96u, status, descriptors, params + 21, scratch + offsets[2], scratch + offsets[17], scratch, scratch + offsets[18], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(57u, status, descriptors, params + 22, scratch + offsets[14], scratch, scratch, scratch + offsets[19], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(57u, status, descriptors, params + 23, scratch + offsets[18], scratch, scratch, scratch + offsets[20], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(24u, status, descriptors, params + 24, scratch + offsets[19], scratch + offsets[20], scratch, scratch + offsets[21], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(7u, status, descriptors, params + 25, scratch + offsets[21], scratch, scratch, scratch + offsets[22], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(49u, status, descriptors, params + 26, scratch + offsets[22], scratch, scratch, scratch + offsets[23], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(7u, status, descriptors, params + 27, scratch + offsets[23], scratch, scratch, scratch + offsets[24], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(146u, status, descriptors, params + 28, scratch + offsets[24], scratch, scratch, pending_6, scratch, temporary);
  if (status->state != 1) return;
  current_6 = pending_6;
}
=== end
=== synthetic_sort_desc#0 singleton#0
function_name: ptir_m2_dc132d209d85d958_r0
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41548 prefix=@runtime tail_bytes=1049
kernel void ptir_m2_dc132d209d85d958_r0(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  ptir_m1_execute(145u, status, descriptors, params + 0, current_0, scratch, scratch, scratch + offsets[0], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== synthetic_sort_desc#0 singleton#1
function_name: ptir_m2_dc132d209d85d958_r1
region: library=1 library_op=2 schedule=3 nodes=1 inputs=1 outputs=2 sinks=0
ok: true
source: bytes=41571 prefix=@runtime tail_bytes=1072
kernel void ptir_m2_dc132d209d85d958_r1(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  ptir_m1_execute(80u, status, descriptors, params + 1, scratch + offsets[0], scratch, scratch, scratch + offsets[1], scratch + offsets[2], temporary);
  if (status->state != 1) return;
}
=== end
=== synthetic_sort_desc#0 singleton#2
function_name: ptir_m2_dc132d209d85d958_r2
region: library=0 library_op=0 schedule=0 nodes=1 inputs=1 outputs=0 sinks=1
ok: true
source: bytes=41573 prefix=@runtime tail_bytes=1074
kernel void ptir_m2_dc132d209d85d958_r2(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  ptir_m1_execute(146u, status, descriptors, params + 2, scratch + offsets[1], scratch, scratch, pending_1, scratch, temporary);
  if (status->state != 1) return;
  current_1 = pending_1;
}
=== end
=== synthetic_sort_desc#0 fused#0
function_name: ptir_m2_dc132d209d85d958_r0
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41548 prefix=@runtime tail_bytes=1049
kernel void ptir_m2_dc132d209d85d958_r0(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  ptir_m1_execute(145u, status, descriptors, params + 0, current_0, scratch, scratch, scratch + offsets[0], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== synthetic_sort_desc#0 fused#1
function_name: ptir_m2_dc132d209d85d958_r1
region: library=1 library_op=2 schedule=3 nodes=1 inputs=1 outputs=2 sinks=0
ok: true
source: bytes=41571 prefix=@runtime tail_bytes=1072
kernel void ptir_m2_dc132d209d85d958_r1(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  ptir_m1_execute(80u, status, descriptors, params + 1, scratch + offsets[0], scratch, scratch, scratch + offsets[1], scratch + offsets[2], temporary);
  if (status->state != 1) return;
}
=== end
=== synthetic_sort_desc#0 fused#2
function_name: ptir_m2_dc132d209d85d958_r2
region: library=0 library_op=0 schedule=0 nodes=2 inputs=2 outputs=0 sinks=2
ok: true
source: bytes=41761 prefix=@runtime tail_bytes=1262
kernel void ptir_m2_dc132d209d85d958_r2(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  ptir_m1_execute(146u, status, descriptors, params + 2, scratch + offsets[1], scratch, scratch, pending_1, scratch, temporary);
  if (status->state != 1) return;
  current_1 = pending_1;
  ptir_m1_execute(146u, status, descriptors, params + 3, scratch + offsets[2], scratch, scratch, pending_2, scratch, temporary);
  if (status->state != 1) return;
  current_2 = pending_2;
}
=== end
=== synthetic_matmul#0 singleton#0
function_name: ptir_m2_92d86f95b487e729_r0
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41548 prefix=@runtime tail_bytes=1049
kernel void ptir_m2_92d86f95b487e729_r0(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  ptir_m1_execute(145u, status, descriptors, params + 0, current_0, scratch, scratch, scratch + offsets[0], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== synthetic_matmul#0 singleton#2
function_name: ptir_m2_92d86f95b487e729_r2
region: library=1 library_op=4 schedule=3 nodes=1 inputs=2 outputs=1 sinks=0
ok: true
source: bytes=41571 prefix=@runtime tail_bytes=1072
kernel void ptir_m2_92d86f95b487e729_r2(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  ptir_m1_execute(85u, status, descriptors, params + 2, scratch + offsets[0], scratch + offsets[1], scratch, scratch + offsets[2], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== synthetic_matmul#0 singleton#3
function_name: ptir_m2_92d86f95b487e729_r3
region: library=0 library_op=0 schedule=0 nodes=1 inputs=1 outputs=0 sinks=1
ok: true
source: bytes=41573 prefix=@runtime tail_bytes=1074
kernel void ptir_m2_92d86f95b487e729_r3(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  ptir_m1_execute(146u, status, descriptors, params + 3, scratch + offsets[2], scratch, scratch, pending_2, scratch, temporary);
  if (status->state != 1) return;
  current_2 = pending_2;
}
=== end
=== synthetic_matmul#0 fused#0
function_name: ptir_m2_92d86f95b487e729_r0
region: library=0 library_op=0 schedule=0 nodes=2 inputs=0 outputs=2 sinks=0
ok: true
source: bytes=41711 prefix=@runtime tail_bytes=1212
kernel void ptir_m2_92d86f95b487e729_r0(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  ptir_m1_execute(145u, status, descriptors, params + 0, current_0, scratch, scratch, scratch + offsets[0], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(145u, status, descriptors, params + 1, current_1, scratch, scratch, scratch + offsets[1], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== synthetic_matmul#0 fused#1
function_name: ptir_m2_92d86f95b487e729_r1
region: library=1 library_op=4 schedule=3 nodes=1 inputs=2 outputs=1 sinks=0
ok: true
source: bytes=41571 prefix=@runtime tail_bytes=1072
kernel void ptir_m2_92d86f95b487e729_r1(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  ptir_m1_execute(85u, status, descriptors, params + 2, scratch + offsets[0], scratch + offsets[1], scratch, scratch + offsets[2], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== synthetic_matmul#0 fused#2
function_name: ptir_m2_92d86f95b487e729_r2
region: library=0 library_op=0 schedule=0 nodes=1 inputs=1 outputs=0 sinks=1
ok: true
source: bytes=41573 prefix=@runtime tail_bytes=1074
kernel void ptir_m2_92d86f95b487e729_r2(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* committed_1 [[buffer(9)]],
    device uchar* pending_1 [[buffer(10)]],
    const device uchar* committed_2 [[buffer(11)]],
    device uchar* pending_2 [[buffer(12)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  const device uchar* current_1 = committed_1;
  const device uchar* current_2 = committed_2;
  ptir_m1_execute(146u, status, descriptors, params + 3, scratch + offsets[2], scratch, scratch, pending_2, scratch, temporary);
  if (status->state != 1) return;
  current_2 = pending_2;
}
=== end
=== synthetic_mtp_drafts#0 singleton#0
function_name: ptir_m2_9eda8aa42400c4c2_r0
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41317 prefix=@runtime tail_bytes=818
kernel void ptir_m2_9eda8aa42400c4c2_r0(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* intrinsic_6 [[buffer(29)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  ptir_m1_execute(160u, status, descriptors, params + 0, intrinsic_6, scratch, scratch, scratch + offsets[0], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== synthetic_mtp_drafts#0 singleton#2
function_name: ptir_m2_9eda8aa42400c4c2_r2
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41273 prefix=@runtime tail_bytes=774
kernel void ptir_m2_9eda8aa42400c4c2_r2(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  ptir_m1_execute(51u, status, descriptors, params + 2, scratch + offsets[1], scratch, scratch, scratch + offsets[2], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== synthetic_mtp_drafts#0 singleton#3
function_name: ptir_m2_9eda8aa42400c4c2_r3
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: true
source: bytes=41272 prefix=@runtime tail_bytes=773
kernel void ptir_m2_9eda8aa42400c4c2_r3(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  ptir_m1_execute(7u, status, descriptors, params + 3, scratch + offsets[2], scratch, scratch, scratch + offsets[3], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== synthetic_mtp_drafts#0 singleton#4
function_name: ptir_m2_9eda8aa42400c4c2_r4
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: true
source: bytes=41286 prefix=@runtime tail_bytes=787
kernel void ptir_m2_9eda8aa42400c4c2_r4(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  ptir_m1_execute(16u, status, descriptors, params + 4, scratch + offsets[0], scratch + offsets[3], scratch, scratch + offsets[4], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== synthetic_mtp_drafts#0 singleton#5
function_name: ptir_m2_9eda8aa42400c4c2_r5
region: library=0 library_op=0 schedule=0 nodes=1 inputs=1 outputs=0 sinks=1
ok: true
source: bytes=41288 prefix=@runtime tail_bytes=789
kernel void ptir_m2_9eda8aa42400c4c2_r5(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  ptir_m1_execute(146u, status, descriptors, params + 5, scratch + offsets[4], scratch, scratch, pending_0, scratch, temporary);
  if (status->state != 1) return;
  current_0 = pending_0;
}
=== end
=== synthetic_mtp_drafts#0 fused#0
function_name: ptir_m2_9eda8aa42400c4c2_r0
region: library=0 library_op=0 schedule=1 nodes=6 inputs=0 outputs=0 sinks=1
ok: true
source: bytes=42253 prefix=@runtime tail_bytes=1754
kernel void ptir_m2_9eda8aa42400c4c2_r0(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    const device uchar* intrinsic_1 [[buffer(30)]],
    const device uchar* intrinsic_6 [[buffer(29)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  ptir_m1_execute(160u, status, descriptors, params + 0, intrinsic_6, scratch, scratch, scratch + offsets[0], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(160u, status, descriptors, params + 1, intrinsic_1, scratch, scratch, scratch + offsets[1], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(51u, status, descriptors, params + 2, scratch + offsets[1], scratch, scratch, scratch + offsets[2], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(7u, status, descriptors, params + 3, scratch + offsets[2], scratch, scratch, scratch + offsets[3], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(16u, status, descriptors, params + 4, scratch + offsets[0], scratch + offsets[3], scratch, scratch + offsets[4], scratch, temporary);
  if (status->state != 1) return;
  ptir_m1_execute(146u, status, descriptors, params + 5, scratch + offsets[4], scratch, scratch, pending_0, scratch, temporary);
  if (status->state != 1) return;
  current_0 = pending_0;
}
=== end
=== synthetic_sink_call#0 singleton#0
function_name: ptir_m2_a297c60f7a5de47e_r0
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41263 prefix=@runtime tail_bytes=764
kernel void ptir_m2_a297c60f7a5de47e_r0(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  ptir_m1_execute(145u, status, descriptors, params + 0, current_0, scratch, scratch, scratch + offsets[0], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== synthetic_sink_call#0 singleton#1
function_name: ptir_m2_a297c60f7a5de47e_r1
region: library=1 library_op=5 schedule=3 nodes=1 inputs=1 outputs=0 sinks=0
ok: true
source: bytes=41261 prefix=@runtime tail_bytes=762
kernel void ptir_m2_a297c60f7a5de47e_r1(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  ptir_m1_execute(162u, status, descriptors, params + 1, scratch + offsets[0], scratch, scratch, scratch, scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== synthetic_sink_call#0 fused#0
function_name: ptir_m2_a297c60f7a5de47e_r0
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: true
source: bytes=41263 prefix=@runtime tail_bytes=764
kernel void ptir_m2_a297c60f7a5de47e_r0(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  ptir_m1_execute(145u, status, descriptors, params + 0, current_0, scratch, scratch, scratch + offsets[0], scratch, temporary);
  if (status->state != 1) return;
}
=== end
=== synthetic_sink_call#0 fused#1
function_name: ptir_m2_a297c60f7a5de47e_r1
region: library=1 library_op=5 schedule=3 nodes=1 inputs=1 outputs=0 sinks=0
ok: true
source: bytes=41261 prefix=@runtime tail_bytes=762
kernel void ptir_m2_a297c60f7a5de47e_r1(
    device M1Status* status [[buffer(0)]],
    const device M1ValueDesc* descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* scratch [[buffer(4)]],
    device uchar* temporary [[buffer(5)]],
    const device uchar* logits [[buffer(6)]],
    const device uchar* committed_0 [[buffer(7)]],
    device uchar* pending_0 [[buffer(8)]],
    uint gid [[thread_position_in_grid]]) {
  if (gid != 0 || status->state != 1) return;
  const device uchar* current_0 = committed_0;
  ptir_m1_execute(162u, status, descriptors, params + 1, scratch + offsets[0], scratch, scratch, scratch, scratch, temporary);
  if (status->state != 1) return;
}
=== end
