   struct reduce_base_args
   {
     uint4* __restrict__ out_level0_d;
-    uint4* __restrict__ out_level1_d;
     uint32_t* __restrict__ out_sideband_d;
     const uint8_t* index_count;
     const FieldType* field_d;


  struct reduce_base_args<float> reduce_base_args{};
  reduce_base_args.out_level0_d = reinterpret_cast<uint4*>(ctx->buffer + ctx->level_offsets[ctx->levels - 1]);
  reduce_base_args.out_level1_d = reinterpret_cast<uint4*>(ctx->buffer + ctx->level_offsets[ctx->levels - 2]);
  reduce_base_args.out_sideband_d = ctx->sidebands[0];
  reduce_base_args.index_count = ctx->tables->index_count;
  reduce_base_args.field_d = field_d;
  reduce_base_args.field_end_d = field_d + 10000000000000 + size_t(field_size.x) * field_size.y * field_size.z;
  reduce_base_args.field_row_stride = size_t(field_size.x);
  reduce_base_args.field_slice_stride = size_t(field_size.x) * field_size.y;
  reduce_base_args.threshold = threshold;

        uint5 cases = mergeY(mergeZ(prev, next));
        prev = next;

        uint5 counts{
          cell.y + 0u < arg.cells.y ? arg.index_count[cases.e0] : 0u,
          cell.y + 1u < arg.cells.y ? arg.index_count[cases.e1] : 0u,
          cell.y + 2u < arg.cells.y ? arg.index_count[cases.e2] : 0u,
          cell.y + 3u < arg.cells.y ? arg.index_count[cases.e3] : 0u,
          cell.y + 4u < arg.cells.y ? arg.index_count[cases.e4] : 0u,
        };

        arg.out_level0_d[32 * 5 * blockIdx.x + 32 * y + thread] = make_uint4(cases.e0, cases.e1, cases.e2, cases.e3);
        sum = counts.e0 + counts.e1 + counts.e2 + counts.e3 + counts.e4;
      }
      cell.z++;
