# oracle dump: pie::metal::pipeline::emit_grouped_topk_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: metal emitter 41->42, and the shared lane record is what moved. The GROUPED (M3) form binds no per-intrinsic buffer -- every rectangle it reads arrives as an address on the lane record -- so `attn_score` was the one id its table refused: the observability slab is the shell's own reservation (engine_metal::scores), not the readout arena, and no row displacement off lane.logits_base reaches it. LaneRecord now carries attn_score_base and attn_score_row_stride (appended, so every field before them keeps its offset; the record is 96 -> 112 bytes), emit_grouped_fused_region gathers the F32 planes from that base, and m3_intrinsic_bindable stops declining the id. What moved in these dumps is exactly the +73 bytes of struct text the two fields add to the M1 and M3 preambles, plus the emitter version spelled into ptir_m3_generic_{ready,commit}_v42. No emitted logic changed for any corpus case: the base corpus reads no score rectangle, so no kernel body here grew an arm. The bump is load-bearing rather than hygienic -- a library cached under 41 reads the 96-byte record, so lane 1 of any grouped table would start sixteen bytes early -- and CUDA bumps 24 -> 25 for the same reason on its own copy of the struct. The reason the record grew at all is the M2 ceiling: the score rectangle costs two of the twelve argument slots the channels grow into, so a ten-channel guest reading scores (trackb-h2o) was refused by a limit that had nothing to do with what it asked for.
=== beam_epilogue#0 singleton#0
function_name: ptir_m3_b4fadd8446eb58dd_r0
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== beam_epilogue#0 singleton#1
function_name: ptir_m3_b4fadd8446eb58dd_r1
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== beam_epilogue#0 singleton#2
function_name: ptir_m3_b4fadd8446eb58dd_r2
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== beam_epilogue#0 singleton#3
function_name: ptir_m3_b4fadd8446eb58dd_r3
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== beam_epilogue#0 singleton#4
function_name: ptir_m3_b4fadd8446eb58dd_r4
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== beam_epilogue#0 singleton#5
function_name: ptir_m3_b4fadd8446eb58dd_r5
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== beam_epilogue#0 singleton#6
function_name: ptir_m3_b4fadd8446eb58dd_r6
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== beam_epilogue#0 singleton#7
function_name: ptir_m3_b4fadd8446eb58dd_r7
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== beam_epilogue#0 singleton#10
function_name: ptir_m3_b4fadd8446eb58dd_r10
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== beam_epilogue#0 singleton#12
function_name: ptir_m3_b4fadd8446eb58dd_r12
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== beam_epilogue#0 singleton#14
function_name: ptir_m3_b4fadd8446eb58dd_r14
region: library=1 library_op=1 schedule=3 nodes=1 inputs=1 outputs=2 sinks=0
ok: true
source: bytes=46797 prefix=@runtime+@grouped tail_bytes=5255

inline uint m3_topk_order_digit(float value, uint pass) {
  if (pass < 8u) {
    if (isnan(value)) return 0u;
    if (value == 0.0f) value = 0.0f;
    const uint bits = as_type<uint>(value);
    const uint ascending =
        (bits & 0x80000000u) != 0u ? ~bits : (bits ^ 0x80000000u);
    return ((~ascending) >> (pass * 4u)) & 15u;
  }
  return isnan(value) ? 1u : 0u;
}

kernel void ptir_m3_b4fadd8446eb58dd_r14(
    const device uchar* lane_bytes [[buffer(0)]],
    const device M1ValueDesc* all_descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* all_scratch [[buffer(4)]],
    const device M3GroupLayout* layout [[buffer(5)]],
    const device uint* channel_bindings [[buffer(6)]],
    device uchar* pending_flags [[buffer(7)]],
    const device uint* lane_indices [[buffer(8)]],
    const device M3RowMeta* all_row_meta [[buffer(9)]],
    const device uint* row_indices [[buffer(10)]],
    uint thread_index [[thread_index_in_threadgroup]],
    uint threads [[threads_per_threadgroup]],
    uint group_position [[threadgroup_position_in_grid]]) {
  (void)channel_bindings;
  (void)pending_flags;
  (void)all_row_meta;
  (void)row_indices;
  if (threads != 256u || layout->reserved1 == 0u) return;
  const uint dispatch_lane = group_position / layout->reserved1;
  const uint row = group_position % layout->reserved1;
  if (dispatch_lane >= layout->lane_count) return;
  const uint lane_index = lane_indices[dispatch_lane];
  const device M3LaneHeader* header =
      reinterpret_cast<const device M3LaneHeader*>(lane_bytes);
  const device M3LaneRecord* lanes =
      reinterpret_cast<const device M3LaneRecord*>(
          lane_bytes + sizeof(M3LaneHeader));
  device M1Status* status =
      reinterpret_cast<device M1Status*>(lanes[lane_index].commit_slot);
  if (status->state != 1u) return;
  const device M1ValueDesc* descriptors =
      all_descriptors + dispatch_lane * layout->value_count;
  const device M1OpParams* lane_params =
      params + dispatch_lane * layout->reserved2;
  device uchar* scratch =
      all_scratch + dispatch_lane * layout->scratch_stride;
  device uchar* temporary = scratch + layout->temporary_offset;
  constexpr uint kInput = 13u;
  constexpr uint kValues = 14u;
  constexpr uint kIndices = 15u;
  constexpr uint k = 2u;

  const M1ValueDesc input_desc = descriptors[kInput];
  if (row >= input_desc.rows) return;
  const uint len = input_desc.last;
  const device float* input =
      reinterpret_cast<const device float*>(scratch + offsets[kInput]) +
      ulong(row) * len;
  device float* top_values =
      reinterpret_cast<device float*>(scratch + offsets[kValues]);
  device uint* top_indices =
      reinterpret_cast<device uint*>(scratch + offsets[kIndices]);
  device uint* order_a =
      reinterpret_cast<device uint*>(
          temporary + ulong(row) * len * 8ul);
  device uint* order_b = order_a + len;
  threadgroup uint digit_offsets[256 * 16];

  for (uint index = thread_index; index < len; index += threads)
    order_a[index] = index;
  threadgroup_barrier(mem_flags::mem_device);
  device uint* input_order = order_a;
  device uint* output_order = order_b;
  const uint chunk_begin =
      uint((ulong(len) * thread_index) / threads);
  const uint chunk_end =
      uint((ulong(len) * (thread_index + 1u)) / threads);
  for (uint pass = 0u; pass < 9u; ++pass) {
    uint digit_counts[16];
    uint digit_written[16];
    for (uint digit = 0u; digit < 16u; ++digit) {
      digit_counts[digit] = 0u;
      digit_written[digit] = 0u;
    }
    for (uint position = chunk_begin; position < chunk_end; ++position) {
      const uint index = input_order[position];
      ++digit_counts[m3_topk_order_digit(input[index], pass)];
    }
    for (uint digit = 0u; digit < 16u; ++digit)
      digit_offsets[thread_index * 16u + digit] = digit_counts[digit];
    threadgroup_barrier(mem_flags::mem_threadgroup);
    if (thread_index == 0u) {
      uint base = 0u;
      for (uint digit = 0u; digit < 16u; ++digit) {
        uint running = base;
        for (uint worker = 0u; worker < threads; ++worker) {
          const uint offset = worker * 16u + digit;
          const uint count_for_worker = digit_offsets[offset];
          digit_offsets[offset] = running;
          running += count_for_worker;
        }
        base = running;
      }
    }
    threadgroup_barrier(mem_flags::mem_threadgroup);
    for (uint position = chunk_begin; position < chunk_end; ++position) {
      const uint index = input_order[position];
      const uint digit = m3_topk_order_digit(input[index], pass);
      output_order[
          digit_offsets[thread_index * 16u + digit] +
          digit_written[digit]++] = index;
    }
    threadgroup_barrier(mem_flags::mem_device);
    device uint* swap = input_order;
    input_order = output_order;
    output_order = swap;
  }
  if (thread_index == 0u) {
    const uint count = min(k, len);
    for (uint position = 0u; position < count; ++position) {
      const uint index = input_order[position];
      top_values[ulong(row) * k + position] = input[index];
      top_indices[ulong(row) * k + position] = index;
    }
  }
  threadgroup_barrier(mem_flags::mem_device);
}
=== end
=== beam_epilogue#0 singleton#15
function_name: ptir_m3_b4fadd8446eb58dd_r15
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== beam_epilogue#0 singleton#16
function_name: ptir_m3_b4fadd8446eb58dd_r16
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== beam_epilogue#0 singleton#18
function_name: ptir_m3_b4fadd8446eb58dd_r18
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== beam_epilogue#0 singleton#25
function_name: ptir_m3_b4fadd8446eb58dd_r25
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== beam_epilogue#0 singleton#26
function_name: ptir_m3_b4fadd8446eb58dd_r26
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== beam_epilogue#0 singleton#28
function_name: ptir_m3_b4fadd8446eb58dd_r28
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== beam_epilogue#0 singleton#30
function_name: ptir_m3_b4fadd8446eb58dd_r30
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== beam_epilogue#0 singleton#31
function_name: ptir_m3_b4fadd8446eb58dd_r31
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== beam_epilogue#0 singleton#35
function_name: ptir_m3_b4fadd8446eb58dd_r35
region: library=0 library_op=0 schedule=1 nodes=1 inputs=3 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== beam_epilogue#0 singleton#43
function_name: ptir_m3_b4fadd8446eb58dd_r43
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== beam_epilogue#0 singleton#48
function_name: ptir_m3_b4fadd8446eb58dd_r48
region: library=0 library_op=0 schedule=0 nodes=1 inputs=1 outputs=0 sinks=1
ok: false
error: invalid grouped TopK library region
=== end
=== beam_epilogue#0 singleton#75
function_name: ptir_m3_b4fadd8446eb58dd_r75
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== beam_epilogue#0 singleton#76
function_name: ptir_m3_b4fadd8446eb58dd_r76
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== beam_epilogue#0 fused#0
function_name: ptir_m3_b4fadd8446eb58dd_r0
region: library=0 library_op=0 schedule=1 nodes=14 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== beam_epilogue#0 fused#1
function_name: ptir_m3_b4fadd8446eb58dd_r1
region: library=1 library_op=1 schedule=3 nodes=1 inputs=1 outputs=2 sinks=0
ok: true
source: bytes=46796 prefix=@runtime+@grouped tail_bytes=5254

inline uint m3_topk_order_digit(float value, uint pass) {
  if (pass < 8u) {
    if (isnan(value)) return 0u;
    if (value == 0.0f) value = 0.0f;
    const uint bits = as_type<uint>(value);
    const uint ascending =
        (bits & 0x80000000u) != 0u ? ~bits : (bits ^ 0x80000000u);
    return ((~ascending) >> (pass * 4u)) & 15u;
  }
  return isnan(value) ? 1u : 0u;
}

kernel void ptir_m3_b4fadd8446eb58dd_r1(
    const device uchar* lane_bytes [[buffer(0)]],
    const device M1ValueDesc* all_descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* all_scratch [[buffer(4)]],
    const device M3GroupLayout* layout [[buffer(5)]],
    const device uint* channel_bindings [[buffer(6)]],
    device uchar* pending_flags [[buffer(7)]],
    const device uint* lane_indices [[buffer(8)]],
    const device M3RowMeta* all_row_meta [[buffer(9)]],
    const device uint* row_indices [[buffer(10)]],
    uint thread_index [[thread_index_in_threadgroup]],
    uint threads [[threads_per_threadgroup]],
    uint group_position [[threadgroup_position_in_grid]]) {
  (void)channel_bindings;
  (void)pending_flags;
  (void)all_row_meta;
  (void)row_indices;
  if (threads != 256u || layout->reserved1 == 0u) return;
  const uint dispatch_lane = group_position / layout->reserved1;
  const uint row = group_position % layout->reserved1;
  if (dispatch_lane >= layout->lane_count) return;
  const uint lane_index = lane_indices[dispatch_lane];
  const device M3LaneHeader* header =
      reinterpret_cast<const device M3LaneHeader*>(lane_bytes);
  const device M3LaneRecord* lanes =
      reinterpret_cast<const device M3LaneRecord*>(
          lane_bytes + sizeof(M3LaneHeader));
  device M1Status* status =
      reinterpret_cast<device M1Status*>(lanes[lane_index].commit_slot);
  if (status->state != 1u) return;
  const device M1ValueDesc* descriptors =
      all_descriptors + dispatch_lane * layout->value_count;
  const device M1OpParams* lane_params =
      params + dispatch_lane * layout->reserved2;
  device uchar* scratch =
      all_scratch + dispatch_lane * layout->scratch_stride;
  device uchar* temporary = scratch + layout->temporary_offset;
  constexpr uint kInput = 13u;
  constexpr uint kValues = 14u;
  constexpr uint kIndices = 15u;
  constexpr uint k = 2u;

  const M1ValueDesc input_desc = descriptors[kInput];
  if (row >= input_desc.rows) return;
  const uint len = input_desc.last;
  const device float* input =
      reinterpret_cast<const device float*>(scratch + offsets[kInput]) +
      ulong(row) * len;
  device float* top_values =
      reinterpret_cast<device float*>(scratch + offsets[kValues]);
  device uint* top_indices =
      reinterpret_cast<device uint*>(scratch + offsets[kIndices]);
  device uint* order_a =
      reinterpret_cast<device uint*>(
          temporary + ulong(row) * len * 8ul);
  device uint* order_b = order_a + len;
  threadgroup uint digit_offsets[256 * 16];

  for (uint index = thread_index; index < len; index += threads)
    order_a[index] = index;
  threadgroup_barrier(mem_flags::mem_device);
  device uint* input_order = order_a;
  device uint* output_order = order_b;
  const uint chunk_begin =
      uint((ulong(len) * thread_index) / threads);
  const uint chunk_end =
      uint((ulong(len) * (thread_index + 1u)) / threads);
  for (uint pass = 0u; pass < 9u; ++pass) {
    uint digit_counts[16];
    uint digit_written[16];
    for (uint digit = 0u; digit < 16u; ++digit) {
      digit_counts[digit] = 0u;
      digit_written[digit] = 0u;
    }
    for (uint position = chunk_begin; position < chunk_end; ++position) {
      const uint index = input_order[position];
      ++digit_counts[m3_topk_order_digit(input[index], pass)];
    }
    for (uint digit = 0u; digit < 16u; ++digit)
      digit_offsets[thread_index * 16u + digit] = digit_counts[digit];
    threadgroup_barrier(mem_flags::mem_threadgroup);
    if (thread_index == 0u) {
      uint base = 0u;
      for (uint digit = 0u; digit < 16u; ++digit) {
        uint running = base;
        for (uint worker = 0u; worker < threads; ++worker) {
          const uint offset = worker * 16u + digit;
          const uint count_for_worker = digit_offsets[offset];
          digit_offsets[offset] = running;
          running += count_for_worker;
        }
        base = running;
      }
    }
    threadgroup_barrier(mem_flags::mem_threadgroup);
    for (uint position = chunk_begin; position < chunk_end; ++position) {
      const uint index = input_order[position];
      const uint digit = m3_topk_order_digit(input[index], pass);
      output_order[
          digit_offsets[thread_index * 16u + digit] +
          digit_written[digit]++] = index;
    }
    threadgroup_barrier(mem_flags::mem_device);
    device uint* swap = input_order;
    input_order = output_order;
    output_order = swap;
  }
  if (thread_index == 0u) {
    const uint count = min(k, len);
    for (uint position = 0u; position < count; ++position) {
      const uint index = input_order[position];
      top_values[ulong(row) * k + position] = input[index];
      top_indices[ulong(row) * k + position] = index;
    }
  }
  threadgroup_barrier(mem_flags::mem_device);
}
=== end
=== beam_epilogue#0 fused#2
function_name: ptir_m3_b4fadd8446eb58dd_r2
region: library=0 library_op=0 schedule=1 nodes=67 inputs=2 outputs=0 sinks=15
ok: false
error: invalid grouped TopK library region
=== end
=== counter_pingpong#0 singleton#0
function_name: ptir_m3_f4e949004f50e90e_r0
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== counter_pingpong#0 singleton#1
function_name: ptir_m3_f4e949004f50e90e_r1
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== counter_pingpong#0 singleton#2
function_name: ptir_m3_f4e949004f50e90e_r2
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== counter_pingpong#0 singleton#3
function_name: ptir_m3_f4e949004f50e90e_r3
region: library=0 library_op=0 schedule=0 nodes=1 inputs=1 outputs=0 sinks=1
ok: false
error: invalid grouped TopK library region
=== end
=== counter_pingpong#0 fused#0
function_name: ptir_m3_f4e949004f50e90e_r0
region: library=0 library_op=0 schedule=1 nodes=5 inputs=0 outputs=0 sinks=2
ok: false
error: invalid grouped TopK library region
=== end
=== dfa_ingraph#0 singleton#0
function_name: ptir_m3_4f4b856e5ef3e8b4_r0
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== dfa_ingraph#0 singleton#1
function_name: ptir_m3_4f4b856e5ef3e8b4_r1
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== dfa_ingraph#0 singleton#3
function_name: ptir_m3_4f4b856e5ef3e8b4_r3
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== dfa_ingraph#0 singleton#4
function_name: ptir_m3_4f4b856e5ef3e8b4_r4
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== dfa_ingraph#0 singleton#5
function_name: ptir_m3_4f4b856e5ef3e8b4_r5
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== dfa_ingraph#0 singleton#6
function_name: ptir_m3_4f4b856e5ef3e8b4_r6
region: library=0 library_op=0 schedule=1 nodes=1 inputs=3 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== dfa_ingraph#0 singleton#7
function_name: ptir_m3_4f4b856e5ef3e8b4_r7
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== dfa_ingraph#0 singleton#8
function_name: ptir_m3_4f4b856e5ef3e8b4_r8
region: library=0 library_op=0 schedule=0 nodes=1 inputs=1 outputs=0 sinks=1
ok: false
error: invalid grouped TopK library region
=== end
=== dfa_ingraph#0 singleton#10
function_name: ptir_m3_4f4b856e5ef3e8b4_r10
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== dfa_ingraph#0 singleton#11
function_name: ptir_m3_4f4b856e5ef3e8b4_r11
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== dfa_ingraph#0 singleton#12
function_name: ptir_m3_4f4b856e5ef3e8b4_r12
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== dfa_ingraph#0 fused#0
function_name: ptir_m3_4f4b856e5ef3e8b4_r0
region: library=0 library_op=0 schedule=1 nodes=15 inputs=0 outputs=0 sinks=2
ok: false
error: invalid grouped TopK library region
=== end
=== extern_contrastive#0 singleton#0
function_name: ptir_m3_d01f407b22863878_r0
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== extern_contrastive#0 singleton#1
function_name: ptir_m3_d01f407b22863878_r1
region: library=0 library_op=0 schedule=0 nodes=1 inputs=1 outputs=0 sinks=1
ok: false
error: invalid grouped TopK library region
=== end
=== extern_contrastive#0 fused#0
function_name: ptir_m3_d01f407b22863878_r0
region: library=0 library_op=0 schedule=1 nodes=2 inputs=0 outputs=0 sinks=1
ok: false
error: invalid grouped TopK library region
=== end
=== greedy_argmax#0 singleton#0
function_name: ptir_m3_f2566cc0c1e0529a_r0
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== greedy_argmax#0 singleton#1
function_name: ptir_m3_f2566cc0c1e0529a_r1
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== greedy_argmax#0 singleton#2
function_name: ptir_m3_f2566cc0c1e0529a_r2
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== greedy_argmax#0 singleton#4
function_name: ptir_m3_f2566cc0c1e0529a_r4
region: library=0 library_op=0 schedule=0 nodes=1 inputs=1 outputs=0 sinks=1
ok: false
error: invalid grouped TopK library region
=== end
=== greedy_argmax#0 fused#0
function_name: ptir_m3_f2566cc0c1e0529a_r0
region: library=0 library_op=0 schedule=1 nodes=6 inputs=0 outputs=0 sinks=2
ok: false
error: invalid grouped TopK library region
=== end
=== lora_prologue#0 singleton#0
function_name: ptir_m3_26849b68dee54237_r0
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== lora_prologue#0 singleton#2
function_name: ptir_m3_26849b68dee54237_r2
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== lora_prologue#0 singleton#3
function_name: ptir_m3_26849b68dee54237_r3
region: library=1 library_op=5 schedule=3 nodes=1 inputs=3 outputs=0 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== lora_prologue#0 fused#0
function_name: ptir_m3_26849b68dee54237_r0
region: library=0 library_op=0 schedule=1 nodes=3 inputs=0 outputs=3 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== lora_prologue#0 fused#1
function_name: ptir_m3_26849b68dee54237_r1
region: library=1 library_op=5 schedule=3 nodes=1 inputs=3 outputs=0 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== matrix_mask_apply_packed#0 singleton#0
function_name: ptir_m3_6612fe00a51d7edc_r0
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== matrix_mask_apply_packed#0 singleton#1
function_name: ptir_m3_6612fe00a51d7edc_r1
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== matrix_mask_apply_packed#0 singleton#2
function_name: ptir_m3_6612fe00a51d7edc_r2
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== matrix_mask_apply_packed#0 singleton#3
function_name: ptir_m3_6612fe00a51d7edc_r3
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== matrix_mask_apply_packed#0 singleton#4
function_name: ptir_m3_6612fe00a51d7edc_r4
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== matrix_mask_apply_packed#0 singleton#5
function_name: ptir_m3_6612fe00a51d7edc_r5
region: library=0 library_op=0 schedule=0 nodes=1 inputs=1 outputs=0 sinks=1
ok: false
error: invalid grouped TopK library region
=== end
=== matrix_mask_apply_packed#0 fused#0
function_name: ptir_m3_6612fe00a51d7edc_r0
region: library=0 library_op=0 schedule=1 nodes=6 inputs=0 outputs=0 sinks=1
ok: false
error: invalid grouped TopK library region
=== end
=== matrix_select_mask#0 singleton#0
function_name: ptir_m3_101abcfb79b08b89_r0
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== matrix_select_mask#0 singleton#1
function_name: ptir_m3_101abcfb79b08b89_r1
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== matrix_select_mask#0 singleton#2
function_name: ptir_m3_101abcfb79b08b89_r2
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== matrix_select_mask#0 singleton#3
function_name: ptir_m3_101abcfb79b08b89_r3
region: library=0 library_op=0 schedule=1 nodes=1 inputs=3 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== matrix_select_mask#0 singleton#4
function_name: ptir_m3_101abcfb79b08b89_r4
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== matrix_select_mask#0 singleton#5
function_name: ptir_m3_101abcfb79b08b89_r5
region: library=0 library_op=0 schedule=0 nodes=1 inputs=1 outputs=0 sinks=1
ok: false
error: invalid grouped TopK library region
=== end
=== matrix_select_mask#0 fused#0
function_name: ptir_m3_101abcfb79b08b89_r0
region: library=0 library_op=0 schedule=1 nodes=6 inputs=0 outputs=0 sinks=1
ok: false
error: invalid grouped TopK library region
=== end
=== mtp_verify_tail#0 singleton#0
function_name: ptir_m3_963137d4ffed0445_r0
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== mtp_verify_tail#0 singleton#2
function_name: ptir_m3_963137d4ffed0445_r2
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== mtp_verify_tail#0 singleton#3
function_name: ptir_m3_963137d4ffed0445_r3
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== mtp_verify_tail#0 singleton#4
function_name: ptir_m3_963137d4ffed0445_r4
region: library=0 library_op=0 schedule=1 nodes=1 inputs=3 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== mtp_verify_tail#0 singleton#5
function_name: ptir_m3_963137d4ffed0445_r5
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== mtp_verify_tail#0 singleton#7
function_name: ptir_m3_963137d4ffed0445_r7
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== mtp_verify_tail#0 singleton#8
function_name: ptir_m3_963137d4ffed0445_r8
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== mtp_verify_tail#0 singleton#9
function_name: ptir_m3_963137d4ffed0445_r9
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== mtp_verify_tail#0 singleton#13
function_name: ptir_m3_963137d4ffed0445_r13
region: library=1 library_op=3 schedule=3 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== mtp_verify_tail#0 singleton#14
function_name: ptir_m3_963137d4ffed0445_r14
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== mtp_verify_tail#0 singleton#15
function_name: ptir_m3_963137d4ffed0445_r15
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== mtp_verify_tail#0 singleton#16
function_name: ptir_m3_963137d4ffed0445_r16
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== mtp_verify_tail#0 singleton#18
function_name: ptir_m3_963137d4ffed0445_r18
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== mtp_verify_tail#0 singleton#21
function_name: ptir_m3_963137d4ffed0445_r21
region: library=0 library_op=0 schedule=0 nodes=1 inputs=1 outputs=0 sinks=1
ok: false
error: invalid grouped TopK library region
=== end
=== mtp_verify_tail#0 fused#0
function_name: ptir_m3_963137d4ffed0445_r0
region: library=0 library_op=0 schedule=1 nodes=13 inputs=0 outputs=3 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== mtp_verify_tail#0 fused#1
function_name: ptir_m3_963137d4ffed0445_r1
region: library=1 library_op=3 schedule=3 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== mtp_verify_tail#0 fused#2
function_name: ptir_m3_963137d4ffed0445_r2
region: library=0 library_op=0 schedule=1 nodes=11 inputs=3 outputs=0 sinks=3
ok: false
error: invalid grouped TopK library region
=== end
=== nucleus_sample#0 singleton#0
function_name: ptir_m3_82921040f7bc2612_r0
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== nucleus_sample#0 singleton#1
function_name: ptir_m3_82921040f7bc2612_r1
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== nucleus_sample#0 singleton#2
function_name: ptir_m3_82921040f7bc2612_r2
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== nucleus_sample#0 singleton#3
function_name: ptir_m3_82921040f7bc2612_r3
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== nucleus_sample#0 singleton#4
function_name: ptir_m3_82921040f7bc2612_r4
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== nucleus_sample#0 singleton#5
function_name: ptir_m3_82921040f7bc2612_r5
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== nucleus_sample#0 singleton#6
function_name: ptir_m3_82921040f7bc2612_r6
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== nucleus_sample#0 singleton#7
function_name: ptir_m3_82921040f7bc2612_r7
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== nucleus_sample#0 singleton#9
function_name: ptir_m3_82921040f7bc2612_r9
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== nucleus_sample#0 singleton#10
function_name: ptir_m3_82921040f7bc2612_r10
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== nucleus_sample#0 singleton#11
function_name: ptir_m3_82921040f7bc2612_r11
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== nucleus_sample#0 singleton#12
function_name: ptir_m3_82921040f7bc2612_r12
region: library=0 library_op=0 schedule=1 nodes=1 inputs=3 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== nucleus_sample#0 singleton#13
function_name: ptir_m3_82921040f7bc2612_r13
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== nucleus_sample#0 singleton#14
function_name: ptir_m3_82921040f7bc2612_r14
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== nucleus_sample#0 singleton#15
function_name: ptir_m3_82921040f7bc2612_r15
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== nucleus_sample#0 singleton#16
function_name: ptir_m3_82921040f7bc2612_r16
region: library=0 library_op=0 schedule=0 nodes=1 inputs=1 outputs=0 sinks=1
ok: false
error: invalid grouped TopK library region
=== end
=== nucleus_sample#0 fused#0
function_name: ptir_m3_82921040f7bc2612_r0
region: library=0 library_op=0 schedule=1 nodes=3 inputs=0 outputs=3 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== nucleus_sample#0 fused#1
function_name: ptir_m3_82921040f7bc2612_r1
region: library=1 library_op=0 schedule=3 nodes=13 inputs=3 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== nucleus_sample#0 fused#2
function_name: ptir_m3_82921040f7bc2612_r2
region: library=0 library_op=0 schedule=0 nodes=1 inputs=1 outputs=0 sinks=1
ok: false
error: invalid grouped TopK library region
=== end
=== pentathlon_iter#0 singleton#0
function_name: ptir_m3_14f667a41eb30669_r0
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== pentathlon_iter#0 singleton#1
function_name: ptir_m3_14f667a41eb30669_r1
region: library=1 library_op=5 schedule=3 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== pentathlon_iter#0 singleton#2
function_name: ptir_m3_14f667a41eb30669_r2
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== pentathlon_iter#0 singleton#3
function_name: ptir_m3_14f667a41eb30669_r3
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== pentathlon_iter#0 singleton#4
function_name: ptir_m3_14f667a41eb30669_r4
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== pentathlon_iter#0 singleton#5
function_name: ptir_m3_14f667a41eb30669_r5
region: library=1 library_op=5 schedule=3 nodes=1 inputs=1 outputs=0 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== pentathlon_iter#0 fused#0
function_name: ptir_m3_14f667a41eb30669_r0
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== pentathlon_iter#0 fused#1
function_name: ptir_m3_14f667a41eb30669_r1
region: library=1 library_op=5 schedule=3 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== pentathlon_iter#0 fused#2
function_name: ptir_m3_14f667a41eb30669_r2
region: library=0 library_op=0 schedule=1 nodes=3 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== pentathlon_iter#0 fused#3
function_name: ptir_m3_14f667a41eb30669_r3
region: library=1 library_op=5 schedule=3 nodes=1 inputs=1 outputs=0 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== pentathlon_iter#1 singleton#0
function_name: ptir_m3_dff3bc6795a9dcdf_r0
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== pentathlon_iter#1 singleton#1
function_name: ptir_m3_dff3bc6795a9dcdf_r1
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== pentathlon_iter#1 singleton#3
function_name: ptir_m3_dff3bc6795a9dcdf_r3
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== pentathlon_iter#1 singleton#4
function_name: ptir_m3_dff3bc6795a9dcdf_r4
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== pentathlon_iter#1 singleton#5
function_name: ptir_m3_dff3bc6795a9dcdf_r5
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== pentathlon_iter#1 singleton#6
function_name: ptir_m3_dff3bc6795a9dcdf_r6
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== pentathlon_iter#1 singleton#7
function_name: ptir_m3_dff3bc6795a9dcdf_r7
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== pentathlon_iter#1 singleton#8
function_name: ptir_m3_dff3bc6795a9dcdf_r8
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== pentathlon_iter#1 singleton#19
function_name: ptir_m3_dff3bc6795a9dcdf_r19
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== pentathlon_iter#1 singleton#20
function_name: ptir_m3_dff3bc6795a9dcdf_r20
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== pentathlon_iter#1 singleton#23
function_name: ptir_m3_dff3bc6795a9dcdf_r23
region: library=0 library_op=0 schedule=1 nodes=1 inputs=3 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== pentathlon_iter#1 singleton#25
function_name: ptir_m3_dff3bc6795a9dcdf_r25
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== pentathlon_iter#1 singleton#28
function_name: ptir_m3_dff3bc6795a9dcdf_r28
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== pentathlon_iter#1 singleton#29
function_name: ptir_m3_dff3bc6795a9dcdf_r29
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== pentathlon_iter#1 singleton#31
function_name: ptir_m3_dff3bc6795a9dcdf_r31
region: library=1 library_op=1 schedule=3 nodes=1 inputs=1 outputs=2 sinks=0
ok: true
source: bytes=46797 prefix=@runtime+@grouped tail_bytes=5255

inline uint m3_topk_order_digit(float value, uint pass) {
  if (pass < 8u) {
    if (isnan(value)) return 0u;
    if (value == 0.0f) value = 0.0f;
    const uint bits = as_type<uint>(value);
    const uint ascending =
        (bits & 0x80000000u) != 0u ? ~bits : (bits ^ 0x80000000u);
    return ((~ascending) >> (pass * 4u)) & 15u;
  }
  return isnan(value) ? 1u : 0u;
}

kernel void ptir_m3_dff3bc6795a9dcdf_r31(
    const device uchar* lane_bytes [[buffer(0)]],
    const device M1ValueDesc* all_descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* all_scratch [[buffer(4)]],
    const device M3GroupLayout* layout [[buffer(5)]],
    const device uint* channel_bindings [[buffer(6)]],
    device uchar* pending_flags [[buffer(7)]],
    const device uint* lane_indices [[buffer(8)]],
    const device M3RowMeta* all_row_meta [[buffer(9)]],
    const device uint* row_indices [[buffer(10)]],
    uint thread_index [[thread_index_in_threadgroup]],
    uint threads [[threads_per_threadgroup]],
    uint group_position [[threadgroup_position_in_grid]]) {
  (void)channel_bindings;
  (void)pending_flags;
  (void)all_row_meta;
  (void)row_indices;
  if (threads != 256u || layout->reserved1 == 0u) return;
  const uint dispatch_lane = group_position / layout->reserved1;
  const uint row = group_position % layout->reserved1;
  if (dispatch_lane >= layout->lane_count) return;
  const uint lane_index = lane_indices[dispatch_lane];
  const device M3LaneHeader* header =
      reinterpret_cast<const device M3LaneHeader*>(lane_bytes);
  const device M3LaneRecord* lanes =
      reinterpret_cast<const device M3LaneRecord*>(
          lane_bytes + sizeof(M3LaneHeader));
  device M1Status* status =
      reinterpret_cast<device M1Status*>(lanes[lane_index].commit_slot);
  if (status->state != 1u) return;
  const device M1ValueDesc* descriptors =
      all_descriptors + dispatch_lane * layout->value_count;
  const device M1OpParams* lane_params =
      params + dispatch_lane * layout->reserved2;
  device uchar* scratch =
      all_scratch + dispatch_lane * layout->scratch_stride;
  device uchar* temporary = scratch + layout->temporary_offset;
  constexpr uint kInput = 30u;
  constexpr uint kValues = 31u;
  constexpr uint kIndices = 32u;
  constexpr uint k = 2u;

  const M1ValueDesc input_desc = descriptors[kInput];
  if (row >= input_desc.rows) return;
  const uint len = input_desc.last;
  const device float* input =
      reinterpret_cast<const device float*>(scratch + offsets[kInput]) +
      ulong(row) * len;
  device float* top_values =
      reinterpret_cast<device float*>(scratch + offsets[kValues]);
  device uint* top_indices =
      reinterpret_cast<device uint*>(scratch + offsets[kIndices]);
  device uint* order_a =
      reinterpret_cast<device uint*>(
          temporary + ulong(row) * len * 8ul);
  device uint* order_b = order_a + len;
  threadgroup uint digit_offsets[256 * 16];

  for (uint index = thread_index; index < len; index += threads)
    order_a[index] = index;
  threadgroup_barrier(mem_flags::mem_device);
  device uint* input_order = order_a;
  device uint* output_order = order_b;
  const uint chunk_begin =
      uint((ulong(len) * thread_index) / threads);
  const uint chunk_end =
      uint((ulong(len) * (thread_index + 1u)) / threads);
  for (uint pass = 0u; pass < 9u; ++pass) {
    uint digit_counts[16];
    uint digit_written[16];
    for (uint digit = 0u; digit < 16u; ++digit) {
      digit_counts[digit] = 0u;
      digit_written[digit] = 0u;
    }
    for (uint position = chunk_begin; position < chunk_end; ++position) {
      const uint index = input_order[position];
      ++digit_counts[m3_topk_order_digit(input[index], pass)];
    }
    for (uint digit = 0u; digit < 16u; ++digit)
      digit_offsets[thread_index * 16u + digit] = digit_counts[digit];
    threadgroup_barrier(mem_flags::mem_threadgroup);
    if (thread_index == 0u) {
      uint base = 0u;
      for (uint digit = 0u; digit < 16u; ++digit) {
        uint running = base;
        for (uint worker = 0u; worker < threads; ++worker) {
          const uint offset = worker * 16u + digit;
          const uint count_for_worker = digit_offsets[offset];
          digit_offsets[offset] = running;
          running += count_for_worker;
        }
        base = running;
      }
    }
    threadgroup_barrier(mem_flags::mem_threadgroup);
    for (uint position = chunk_begin; position < chunk_end; ++position) {
      const uint index = input_order[position];
      const uint digit = m3_topk_order_digit(input[index], pass);
      output_order[
          digit_offsets[thread_index * 16u + digit] +
          digit_written[digit]++] = index;
    }
    threadgroup_barrier(mem_flags::mem_device);
    device uint* swap = input_order;
    input_order = output_order;
    output_order = swap;
  }
  if (thread_index == 0u) {
    const uint count = min(k, len);
    for (uint position = 0u; position < count; ++position) {
      const uint index = input_order[position];
      top_values[ulong(row) * k + position] = input[index];
      top_indices[ulong(row) * k + position] = index;
    }
  }
  threadgroup_barrier(mem_flags::mem_device);
}
=== end
=== pentathlon_iter#1 singleton#34
function_name: ptir_m3_dff3bc6795a9dcdf_r34
region: library=0 library_op=0 schedule=0 nodes=1 inputs=1 outputs=0 sinks=1
ok: false
error: invalid grouped TopK library region
=== end
=== pentathlon_iter#1 fused#0
function_name: ptir_m3_dff3bc6795a9dcdf_r0
region: library=0 library_op=0 schedule=1 nodes=31 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== pentathlon_iter#1 fused#1
function_name: ptir_m3_dff3bc6795a9dcdf_r1
region: library=1 library_op=1 schedule=3 nodes=1 inputs=1 outputs=2 sinks=0
ok: true
source: bytes=46796 prefix=@runtime+@grouped tail_bytes=5254

inline uint m3_topk_order_digit(float value, uint pass) {
  if (pass < 8u) {
    if (isnan(value)) return 0u;
    if (value == 0.0f) value = 0.0f;
    const uint bits = as_type<uint>(value);
    const uint ascending =
        (bits & 0x80000000u) != 0u ? ~bits : (bits ^ 0x80000000u);
    return ((~ascending) >> (pass * 4u)) & 15u;
  }
  return isnan(value) ? 1u : 0u;
}

kernel void ptir_m3_dff3bc6795a9dcdf_r1(
    const device uchar* lane_bytes [[buffer(0)]],
    const device M1ValueDesc* all_descriptors [[buffer(1)]],
    const device M1OpParams* params [[buffer(2)]],
    const device uint* offsets [[buffer(3)]],
    device uchar* all_scratch [[buffer(4)]],
    const device M3GroupLayout* layout [[buffer(5)]],
    const device uint* channel_bindings [[buffer(6)]],
    device uchar* pending_flags [[buffer(7)]],
    const device uint* lane_indices [[buffer(8)]],
    const device M3RowMeta* all_row_meta [[buffer(9)]],
    const device uint* row_indices [[buffer(10)]],
    uint thread_index [[thread_index_in_threadgroup]],
    uint threads [[threads_per_threadgroup]],
    uint group_position [[threadgroup_position_in_grid]]) {
  (void)channel_bindings;
  (void)pending_flags;
  (void)all_row_meta;
  (void)row_indices;
  if (threads != 256u || layout->reserved1 == 0u) return;
  const uint dispatch_lane = group_position / layout->reserved1;
  const uint row = group_position % layout->reserved1;
  if (dispatch_lane >= layout->lane_count) return;
  const uint lane_index = lane_indices[dispatch_lane];
  const device M3LaneHeader* header =
      reinterpret_cast<const device M3LaneHeader*>(lane_bytes);
  const device M3LaneRecord* lanes =
      reinterpret_cast<const device M3LaneRecord*>(
          lane_bytes + sizeof(M3LaneHeader));
  device M1Status* status =
      reinterpret_cast<device M1Status*>(lanes[lane_index].commit_slot);
  if (status->state != 1u) return;
  const device M1ValueDesc* descriptors =
      all_descriptors + dispatch_lane * layout->value_count;
  const device M1OpParams* lane_params =
      params + dispatch_lane * layout->reserved2;
  device uchar* scratch =
      all_scratch + dispatch_lane * layout->scratch_stride;
  device uchar* temporary = scratch + layout->temporary_offset;
  constexpr uint kInput = 30u;
  constexpr uint kValues = 31u;
  constexpr uint kIndices = 32u;
  constexpr uint k = 2u;

  const M1ValueDesc input_desc = descriptors[kInput];
  if (row >= input_desc.rows) return;
  const uint len = input_desc.last;
  const device float* input =
      reinterpret_cast<const device float*>(scratch + offsets[kInput]) +
      ulong(row) * len;
  device float* top_values =
      reinterpret_cast<device float*>(scratch + offsets[kValues]);
  device uint* top_indices =
      reinterpret_cast<device uint*>(scratch + offsets[kIndices]);
  device uint* order_a =
      reinterpret_cast<device uint*>(
          temporary + ulong(row) * len * 8ul);
  device uint* order_b = order_a + len;
  threadgroup uint digit_offsets[256 * 16];

  for (uint index = thread_index; index < len; index += threads)
    order_a[index] = index;
  threadgroup_barrier(mem_flags::mem_device);
  device uint* input_order = order_a;
  device uint* output_order = order_b;
  const uint chunk_begin =
      uint((ulong(len) * thread_index) / threads);
  const uint chunk_end =
      uint((ulong(len) * (thread_index + 1u)) / threads);
  for (uint pass = 0u; pass < 9u; ++pass) {
    uint digit_counts[16];
    uint digit_written[16];
    for (uint digit = 0u; digit < 16u; ++digit) {
      digit_counts[digit] = 0u;
      digit_written[digit] = 0u;
    }
    for (uint position = chunk_begin; position < chunk_end; ++position) {
      const uint index = input_order[position];
      ++digit_counts[m3_topk_order_digit(input[index], pass)];
    }
    for (uint digit = 0u; digit < 16u; ++digit)
      digit_offsets[thread_index * 16u + digit] = digit_counts[digit];
    threadgroup_barrier(mem_flags::mem_threadgroup);
    if (thread_index == 0u) {
      uint base = 0u;
      for (uint digit = 0u; digit < 16u; ++digit) {
        uint running = base;
        for (uint worker = 0u; worker < threads; ++worker) {
          const uint offset = worker * 16u + digit;
          const uint count_for_worker = digit_offsets[offset];
          digit_offsets[offset] = running;
          running += count_for_worker;
        }
        base = running;
      }
    }
    threadgroup_barrier(mem_flags::mem_threadgroup);
    for (uint position = chunk_begin; position < chunk_end; ++position) {
      const uint index = input_order[position];
      const uint digit = m3_topk_order_digit(input[index], pass);
      output_order[
          digit_offsets[thread_index * 16u + digit] +
          digit_written[digit]++] = index;
    }
    threadgroup_barrier(mem_flags::mem_device);
    device uint* swap = input_order;
    input_order = output_order;
    output_order = swap;
  }
  if (thread_index == 0u) {
    const uint count = min(k, len);
    for (uint position = 0u; position < count; ++position) {
      const uint index = input_order[position];
      top_values[ulong(row) * k + position] = input[index];
      top_indices[ulong(row) * k + position] = index;
    }
  }
  threadgroup_barrier(mem_flags::mem_device);
}
=== end
=== pentathlon_iter#1 fused#2
function_name: ptir_m3_dff3bc6795a9dcdf_r2
region: library=0 library_op=0 schedule=1 nodes=6 inputs=2 outputs=0 sinks=3
ok: false
error: invalid grouped TopK library region
=== end
=== pivot_predicates_multistage#1 singleton#0
function_name: ptir_m3_09a2719886970a85_r0
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== pivot_predicates_multistage#1 singleton#1
function_name: ptir_m3_09a2719886970a85_r1
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== pivot_predicates_multistage#1 singleton#2
function_name: ptir_m3_09a2719886970a85_r2
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== pivot_predicates_multistage#1 singleton#3
function_name: ptir_m3_09a2719886970a85_r3
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== pivot_predicates_multistage#1 singleton#4
function_name: ptir_m3_09a2719886970a85_r4
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== pivot_predicates_multistage#1 singleton#5
function_name: ptir_m3_09a2719886970a85_r5
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== pivot_predicates_multistage#1 singleton#7
function_name: ptir_m3_09a2719886970a85_r7
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== pivot_predicates_multistage#1 singleton#8
function_name: ptir_m3_09a2719886970a85_r8
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== pivot_predicates_multistage#1 singleton#10
function_name: ptir_m3_09a2719886970a85_r10
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== pivot_predicates_multistage#1 singleton#12
function_name: ptir_m3_09a2719886970a85_r12
region: library=0 library_op=0 schedule=0 nodes=1 inputs=1 outputs=0 sinks=1
ok: false
error: invalid grouped TopK library region
=== end
=== pivot_predicates_multistage#1 fused#0
function_name: ptir_m3_09a2719886970a85_r0
region: library=0 library_op=0 schedule=1 nodes=14 inputs=0 outputs=0 sinks=2
ok: false
error: invalid grouped TopK library region
=== end
=== section3_masked_gumbel#0 singleton#0
function_name: ptir_m3_4cf3882bf2fff913_r0
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== section3_masked_gumbel#0 singleton#1
function_name: ptir_m3_4cf3882bf2fff913_r1
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== section3_masked_gumbel#0 singleton#2
function_name: ptir_m3_4cf3882bf2fff913_r2
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== section3_masked_gumbel#0 singleton#4
function_name: ptir_m3_4cf3882bf2fff913_r4
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== section3_masked_gumbel#0 singleton#5
function_name: ptir_m3_4cf3882bf2fff913_r5
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== section3_masked_gumbel#0 singleton#6
function_name: ptir_m3_4cf3882bf2fff913_r6
region: library=0 library_op=0 schedule=1 nodes=1 inputs=3 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== section3_masked_gumbel#0 singleton#7
function_name: ptir_m3_4cf3882bf2fff913_r7
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== section3_masked_gumbel#0 singleton#8
function_name: ptir_m3_4cf3882bf2fff913_r8
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== section3_masked_gumbel#0 singleton#9
function_name: ptir_m3_4cf3882bf2fff913_r9
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== section3_masked_gumbel#0 singleton#11
function_name: ptir_m3_4cf3882bf2fff913_r11
region: library=0 library_op=0 schedule=0 nodes=1 inputs=1 outputs=0 sinks=1
ok: false
error: invalid grouped TopK library region
=== end
=== section3_masked_gumbel#0 fused#0
function_name: ptir_m3_4cf3882bf2fff913_r0
region: library=0 library_op=0 schedule=1 nodes=19 inputs=0 outputs=0 sinks=4
ok: false
error: invalid grouped TopK library region
=== end
=== structured_masks#0 singleton#0
function_name: ptir_m3_1752e0dfb0bf5827_r0
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== structured_masks#0 singleton#3
function_name: ptir_m3_1752e0dfb0bf5827_r3
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== structured_masks#0 singleton#4
function_name: ptir_m3_1752e0dfb0bf5827_r4
region: library=0 library_op=0 schedule=0 nodes=1 inputs=1 outputs=0 sinks=1
ok: false
error: invalid grouped TopK library region
=== end
=== structured_masks#0 singleton#5
function_name: ptir_m3_1752e0dfb0bf5827_r5
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== structured_masks#0 singleton#7
function_name: ptir_m3_1752e0dfb0bf5827_r7
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== structured_masks#0 singleton#9
function_name: ptir_m3_1752e0dfb0bf5827_r9
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== structured_masks#0 singleton#10
function_name: ptir_m3_1752e0dfb0bf5827_r10
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== structured_masks#0 singleton#11
function_name: ptir_m3_1752e0dfb0bf5827_r11
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== structured_masks#0 singleton#12
function_name: ptir_m3_1752e0dfb0bf5827_r12
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== structured_masks#0 singleton#14
function_name: ptir_m3_1752e0dfb0bf5827_r14
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== structured_masks#0 singleton#15
function_name: ptir_m3_1752e0dfb0bf5827_r15
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== structured_masks#0 singleton#16
function_name: ptir_m3_1752e0dfb0bf5827_r16
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== structured_masks#0 singleton#17
function_name: ptir_m3_1752e0dfb0bf5827_r17
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== structured_masks#0 singleton#24
function_name: ptir_m3_1752e0dfb0bf5827_r24
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== structured_masks#0 singleton#25
function_name: ptir_m3_1752e0dfb0bf5827_r25
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== structured_masks#0 singleton#26
function_name: ptir_m3_1752e0dfb0bf5827_r26
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== structured_masks#0 fused#0
function_name: ptir_m3_1752e0dfb0bf5827_r0
region: library=0 library_op=0 schedule=1 nodes=29 inputs=0 outputs=0 sinks=4
ok: false
error: invalid grouped TopK library region
=== end
=== synthetic_sort_desc#0 singleton#0
function_name: ptir_m3_dc132d209d85d958_r0
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== synthetic_sort_desc#0 singleton#1
function_name: ptir_m3_dc132d209d85d958_r1
region: library=1 library_op=2 schedule=3 nodes=1 inputs=1 outputs=2 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== synthetic_sort_desc#0 singleton#2
function_name: ptir_m3_dc132d209d85d958_r2
region: library=0 library_op=0 schedule=0 nodes=1 inputs=1 outputs=0 sinks=1
ok: false
error: invalid grouped TopK library region
=== end
=== synthetic_sort_desc#0 fused#0
function_name: ptir_m3_dc132d209d85d958_r0
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== synthetic_sort_desc#0 fused#1
function_name: ptir_m3_dc132d209d85d958_r1
region: library=1 library_op=2 schedule=3 nodes=1 inputs=1 outputs=2 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== synthetic_sort_desc#0 fused#2
function_name: ptir_m3_dc132d209d85d958_r2
region: library=0 library_op=0 schedule=0 nodes=2 inputs=2 outputs=0 sinks=2
ok: false
error: invalid grouped TopK library region
=== end
=== synthetic_matmul#0 singleton#0
function_name: ptir_m3_92d86f95b487e729_r0
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== synthetic_matmul#0 singleton#2
function_name: ptir_m3_92d86f95b487e729_r2
region: library=1 library_op=4 schedule=3 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== synthetic_matmul#0 singleton#3
function_name: ptir_m3_92d86f95b487e729_r3
region: library=0 library_op=0 schedule=0 nodes=1 inputs=1 outputs=0 sinks=1
ok: false
error: invalid grouped TopK library region
=== end
=== synthetic_matmul#0 fused#0
function_name: ptir_m3_92d86f95b487e729_r0
region: library=0 library_op=0 schedule=0 nodes=2 inputs=0 outputs=2 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== synthetic_matmul#0 fused#1
function_name: ptir_m3_92d86f95b487e729_r1
region: library=1 library_op=4 schedule=3 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== synthetic_matmul#0 fused#2
function_name: ptir_m3_92d86f95b487e729_r2
region: library=0 library_op=0 schedule=0 nodes=1 inputs=1 outputs=0 sinks=1
ok: false
error: invalid grouped TopK library region
=== end
=== synthetic_mtp_drafts#0 singleton#0
function_name: ptir_m3_9eda8aa42400c4c2_r0
region: library=0 library_op=0 schedule=1 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== synthetic_mtp_drafts#0 singleton#2
function_name: ptir_m3_9eda8aa42400c4c2_r2
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== synthetic_mtp_drafts#0 singleton#3
function_name: ptir_m3_9eda8aa42400c4c2_r3
region: library=0 library_op=0 schedule=1 nodes=1 inputs=1 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== synthetic_mtp_drafts#0 singleton#4
function_name: ptir_m3_9eda8aa42400c4c2_r4
region: library=0 library_op=0 schedule=1 nodes=1 inputs=2 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== synthetic_mtp_drafts#0 singleton#5
function_name: ptir_m3_9eda8aa42400c4c2_r5
region: library=0 library_op=0 schedule=0 nodes=1 inputs=1 outputs=0 sinks=1
ok: false
error: invalid grouped TopK library region
=== end
=== synthetic_mtp_drafts#0 fused#0
function_name: ptir_m3_9eda8aa42400c4c2_r0
region: library=0 library_op=0 schedule=1 nodes=6 inputs=0 outputs=0 sinks=1
ok: false
error: invalid grouped TopK library region
=== end
=== synthetic_sink_call#0 singleton#0
function_name: ptir_m3_a297c60f7a5de47e_r0
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== synthetic_sink_call#0 singleton#1
function_name: ptir_m3_a297c60f7a5de47e_r1
region: library=1 library_op=5 schedule=3 nodes=1 inputs=1 outputs=0 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== synthetic_sink_call#0 fused#0
function_name: ptir_m3_a297c60f7a5de47e_r0
region: library=0 library_op=0 schedule=0 nodes=1 inputs=0 outputs=1 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
=== synthetic_sink_call#0 fused#1
function_name: ptir_m3_a297c60f7a5de47e_r1
region: library=1 library_op=5 schedule=3 nodes=1 inputs=1 outputs=0 sinks=0
ok: false
error: invalid grouped TopK library region
=== end
