// language: metal3.0
#include <metal_stdlib>
#include <simd/simd.h>

using metal::uint;
struct DefaultConstructible {
    template<typename T>
    operator T() && {
        return T {};
    }
};

struct _mslBufferSizes {
    uint size1;
    uint size2;
    uint size3;
    uint size4;
    uint size5;
    uint size6;
    uint size8;
};

struct Config {
    uint width_in_tiles;
    uint height_in_tiles;
    uint target_width;
    uint target_height;
    uint base_color;
    uint n_drawobj;
    uint n_path;
    uint n_clip;
    uint bin_data_start;
    uint pathtag_base;
    uint pathdata_base;
    uint drawtag_base;
    uint drawdata_base;
    uint transform_base;
    uint style_base;
    uint lines_size;
    uint binning_size;
    uint tiles_size;
    uint seg_counts_size;
    uint segments_size;
    uint blend_size;
    uint ptcl_size;
};
struct BumpAllocators {
    metal::atomic_uint failed;
    metal::atomic_uint binning;
    metal::atomic_uint ptcl;
    metal::atomic_uint tile;
    metal::atomic_uint seg_counts;
    metal::atomic_uint segments;
    metal::atomic_uint blend;
    metal::atomic_uint lines;
};
struct IndirectCount {
    uint count_x;
    uint count_y;
    uint count_z;
};
struct DrawMonoid {
    uint path_ix;
    uint clip_ix;
    uint scene_offset;
    uint info_offset;
};
struct CmdFill {
    uint size_and_rule;
    uint seg_data;
    int backdrop;
};
struct CmdStroke {
    uint tile;
    float half_width;
};
struct CmdJump {
    uint new_ix;
};
struct CmdColor {
    uint rgba_color;
};
struct CmdLinGrad {
    uint index;
    uint extend_mode;
    float line_x;
    float line_y;
    float line_c;
};
struct CmdRadGrad {
    uint index;
    uint extend_mode;
    char _pad2[8];
    metal::float4 matrx;
    metal::float2 xlat;
    float focal_x;
    float radius;
    uint kind;
    uint flags;
};
struct CmdSweepGrad {
    uint index;
    uint extend_mode;
    char _pad2[8];
    metal::float4 matrx;
    metal::float2 xlat;
    float t0_;
    float t1_;
};
struct CmdImage {
    metal::float4 matrx;
    metal::float2 xlat;
    metal::float2 atlas_offset;
    metal::float2 extents;
};
struct CmdEndClip {
    uint blend;
    float alpha;
};
struct Path {
    metal::uint4 bbox;
    uint tiles;
};
struct Tile {
    int backdrop;
    uint segment_count_or_ix;
};
typedef uint type_7[1];
typedef DrawMonoid type_8[1];
struct BinHeader {
    uint element_count;
    uint chunk_offset;
};
typedef BinHeader type_9[1];
typedef Path type_10[1];
typedef Tile type_11[1];
struct type_12 {
    metal::atomic_uint inner[256];
};
struct type_13 {
    type_12 inner[8];
};
struct type_14 {
    uint inner[256];
};
constant uint TILE_WIDTH = 16u;
constant uint TILE_HEIGHT = 16u;
constant uint N_TILE_X = 16u;
constant uint N_TILE_Y = 16u;
constant uint N_TILE = 256u;
constant float TILE_SCALE = 0.0625;
constant uint BLEND_STACK_SPLIT = 4u;
constant uint RAD_GRAD_KIND_CIRCULAR = 1u;
constant uint RAD_GRAD_KIND_STRIP = 2u;
constant uint RAD_GRAD_KIND_FOCAL_ON_CIRCLE = 3u;
constant uint RAD_GRAD_KIND_CONE = 4u;
constant uint RAD_GRAD_SWAPPED = 1u;
constant uint STAGE_BINNING = 1u;
constant uint STAGE_TILE_ALLOC = 2u;
constant uint STAGE_FLATTEN = 4u;
constant uint STAGE_PATH_COUNT = 8u;
constant uint STAGE_COARSE = 16u;
constant uint DRAWTAG_NOP = 0u;
constant uint DRAWTAG_FILL_COLOR = 68u;
constant uint DRAWTAG_FILL_LIN_GRADIENT = 276u;
constant uint DRAWTAG_FILL_RAD_GRADIENT = 668u;
constant uint DRAWTAG_FILL_SWEEP_GRADIENT = 596u;
constant uint DRAWTAG_FILL_IMAGE = 584u;
constant uint DRAWTAG_BEGIN_CLIP = 9u;
constant uint DRAWTAG_END_CLIP = 33u;
constant uint DRAW_INFO_FLAGS_FILL_RULE_BIT = 1u;
constant uint PTCL_INITIAL_ALLOC = 64u;
constant uint PTCL_INCREMENT = 256u;
constant uint PTCL_HEADROOM = 2u;
constant uint CMD_END = 0u;
constant uint CMD_FILL = 1u;
constant uint CMD_STROKE = 2u;
constant uint CMD_SOLID = 3u;
constant uint CMD_COLOR = 5u;
constant uint CMD_LIN_GRAD = 6u;
constant uint CMD_RAD_GRAD = 7u;
constant uint CMD_SWEEP_GRAD = 8u;
constant uint CMD_IMAGE = 9u;
constant uint CMD_BEGIN_CLIP = 10u;
constant uint CMD_END_CLIP = 11u;
constant uint CMD_JUMP = 12u;
constant uint WG_SIZE = 256u;
constant uint N_SLICE = 8u;

DrawMonoid draw_monoid_identity(
) {
    return DrawMonoid {};
}

DrawMonoid combine_draw_monoid(
    DrawMonoid a,
    DrawMonoid b
) {
    DrawMonoid c = {};
    c.path_ix = a.path_ix + b.path_ix;
    c.clip_ix = a.clip_ix + b.clip_ix;
    c.scene_offset = a.scene_offset + b.scene_offset;
    c.info_offset = a.info_offset + b.info_offset;
    DrawMonoid _e19 = c;
    return _e19;
}

DrawMonoid map_draw_tag(
    uint tag_word
) {
    DrawMonoid c_1 = {};
    c_1.path_ix = static_cast<uint>(tag_word != DRAWTAG_NOP);
    c_1.clip_ix = tag_word & 1u;
    c_1.scene_offset = (tag_word >> 2u) & 7u;
    c_1.info_offset = (tag_word >> 6u) & 15u;
    DrawMonoid _e19 = c_1;
    return _e19;
}

void alloc_cmd(
    uint size,
    constant Config& config,
    device BumpAllocators& bump,
    device type_7& ptcl,
    thread uint& cmd_offset,
    thread uint& cmd_limit,
    constant _mslBufferSizes& _buffer_sizes
) {
    uint new_cmd = {};
    uint _e2 = cmd_offset;
    uint _e5 = cmd_limit;
    if ((_e2 + size) >= _e5) {
        uint _e9 = config.width_in_tiles;
        uint _e12 = config.height_in_tiles;
        uint ptcl_dyn_start = (_e9 * _e12) * PTCL_INITIAL_ALLOC;
        uint _e19 = metal::atomic_fetch_add_explicit(&bump.ptcl, PTCL_INCREMENT, metal::memory_order_relaxed);
        new_cmd = ptcl_dyn_start + _e19;
        uint _e22 = new_cmd;
        uint _e27 = config.ptcl_size;
        if ((_e22 + PTCL_INCREMENT) > _e27) {
            new_cmd = 0u;
            uint _e33 = metal::atomic_fetch_or_explicit(&bump.failed, STAGE_COARSE, metal::memory_order_relaxed);
        }
        uint _e36 = cmd_offset;
        if (uint(_e36) < 1 + (_buffer_sizes.size8 - 0 - 4) / 4) {
            ptcl[_e36] = CMD_JUMP;
        }
        uint _e41 = cmd_offset;
        uint _e43 = _e41 + 1u;
        uint _e45 = new_cmd;
        if (uint(_e43) < 1 + (_buffer_sizes.size8 - 0 - 4) / 4) {
            ptcl[_e43] = _e45;
        }
        uint _e47 = new_cmd;
        cmd_offset = _e47;
        uint _e50 = cmd_offset;
        cmd_limit = _e50 + 254u;
        return;
    } else {
        return;
    }
}

void write_path(
    Tile tile,
    uint tile_ix,
    uint draw_flags,
    constant Config& config,
    device type_11& tiles,
    device BumpAllocators& bump,
    device type_7& ptcl,
    thread uint& cmd_offset,
    thread uint& cmd_limit,
    constant _mslBufferSizes& _buffer_sizes
) {
    uint seg_ix = {};
    uint n_segs = tile.segment_count_or_ix;
    if (n_segs != 0u) {
        uint _e8 = metal::atomic_fetch_add_explicit(&bump.segments, n_segs, metal::memory_order_relaxed);
        seg_ix = _e8;
        uint _e13 = seg_ix;
        if (uint(tile_ix) < 1 + (_buffer_sizes.size6 - 0 - 8) / 8) {
            tiles[tile_ix].segment_count_or_ix = ~(_e13);
        }
        alloc_cmd(4u, config, bump, ptcl, cmd_offset, cmd_limit, _buffer_sizes);
        uint _e18 = cmd_offset;
        if (uint(_e18) < 1 + (_buffer_sizes.size8 - 0 - 4) / 4) {
            ptcl[_e18] = CMD_FILL;
        }
        bool even_odd = (draw_flags & DRAW_INFO_FLAGS_FILL_RULE_BIT) != 0u;
        uint size_and_rule = (n_segs << 1u) | static_cast<uint>(even_odd);
        uint _e29 = seg_ix;
        CmdFill fill = CmdFill {size_and_rule, _e29, tile.backdrop};
        uint _e34 = cmd_offset;
        uint _e36 = _e34 + 1u;
        if (uint(_e36) < 1 + (_buffer_sizes.size8 - 0 - 4) / 4) {
            ptcl[_e36] = fill.size_and_rule;
        }
        uint _e41 = cmd_offset;
        uint _e43 = _e41 + 2u;
        if (uint(_e43) < 1 + (_buffer_sizes.size8 - 0 - 4) / 4) {
            ptcl[_e43] = fill.seg_data;
        }
        uint _e48 = cmd_offset;
        uint _e50 = _e48 + 3u;
        if (uint(_e50) < 1 + (_buffer_sizes.size8 - 0 - 4) / 4) {
            ptcl[_e50] = static_cast<uint>(fill.backdrop);
        }
        uint _e56 = cmd_offset;
        cmd_offset = _e56 + 4u;
        return;
    } else {
        alloc_cmd(1u, config, bump, ptcl, cmd_offset, cmd_limit, _buffer_sizes);
        uint _e61 = cmd_offset;
        if (uint(_e61) < 1 + (_buffer_sizes.size8 - 0 - 4) / 4) {
            ptcl[_e61] = CMD_SOLID;
        }
        uint _e66 = cmd_offset;
        cmd_offset = _e66 + 1u;
        return;
    }
}

void write_color(
    CmdColor color,
    constant Config& config,
    device BumpAllocators& bump,
    device type_7& ptcl,
    thread uint& cmd_offset,
    thread uint& cmd_limit,
    constant _mslBufferSizes& _buffer_sizes
) {
    alloc_cmd(2u, config, bump, ptcl, cmd_offset, cmd_limit, _buffer_sizes);
    uint _e4 = cmd_offset;
    if (uint(_e4) < 1 + (_buffer_sizes.size8 - 0 - 4) / 4) {
        ptcl[_e4] = CMD_COLOR;
    }
    uint _e9 = cmd_offset;
    uint _e11 = _e9 + 1u;
    if (uint(_e11) < 1 + (_buffer_sizes.size8 - 0 - 4) / 4) {
        ptcl[_e11] = color.rgba_color;
    }
    uint _e16 = cmd_offset;
    cmd_offset = _e16 + 2u;
    return;
}

void write_grad(
    uint ty,
    uint index,
    uint info_offset,
    constant Config& config,
    device BumpAllocators& bump,
    device type_7& ptcl,
    thread uint& cmd_offset,
    thread uint& cmd_limit,
    constant _mslBufferSizes& _buffer_sizes
) {
    alloc_cmd(3u, config, bump, ptcl, cmd_offset, cmd_limit, _buffer_sizes);
    uint _e6 = cmd_offset;
    if (uint(_e6) < 1 + (_buffer_sizes.size8 - 0 - 4) / 4) {
        ptcl[_e6] = ty;
    }
    uint _e10 = cmd_offset;
    uint _e12 = _e10 + 1u;
    if (uint(_e12) < 1 + (_buffer_sizes.size8 - 0 - 4) / 4) {
        ptcl[_e12] = index;
    }
    uint _e16 = cmd_offset;
    uint _e18 = _e16 + 2u;
    if (uint(_e18) < 1 + (_buffer_sizes.size8 - 0 - 4) / 4) {
        ptcl[_e18] = info_offset;
    }
    uint _e22 = cmd_offset;
    cmd_offset = _e22 + 3u;
    return;
}

void write_image(
    uint info_offset_1,
    constant Config& config,
    device BumpAllocators& bump,
    device type_7& ptcl,
    thread uint& cmd_offset,
    thread uint& cmd_limit,
    constant _mslBufferSizes& _buffer_sizes
) {
    alloc_cmd(2u, config, bump, ptcl, cmd_offset, cmd_limit, _buffer_sizes);
    uint _e4 = cmd_offset;
    if (uint(_e4) < 1 + (_buffer_sizes.size8 - 0 - 4) / 4) {
        ptcl[_e4] = CMD_IMAGE;
    }
    uint _e9 = cmd_offset;
    uint _e11 = _e9 + 1u;
    if (uint(_e11) < 1 + (_buffer_sizes.size8 - 0 - 4) / 4) {
        ptcl[_e11] = info_offset_1;
    }
    uint _e15 = cmd_offset;
    cmd_offset = _e15 + 2u;
    return;
}

void write_begin_clip(
    constant Config& config,
    device BumpAllocators& bump,
    device type_7& ptcl,
    thread uint& cmd_offset,
    thread uint& cmd_limit,
    constant _mslBufferSizes& _buffer_sizes
) {
    alloc_cmd(1u, config, bump, ptcl, cmd_offset, cmd_limit, _buffer_sizes);
    uint _e3 = cmd_offset;
    if (uint(_e3) < 1 + (_buffer_sizes.size8 - 0 - 4) / 4) {
        ptcl[_e3] = CMD_BEGIN_CLIP;
    }
    uint _e8 = cmd_offset;
    cmd_offset = _e8 + 1u;
    return;
}

void write_end_clip(
    CmdEndClip end_clip,
    constant Config& config,
    device BumpAllocators& bump,
    device type_7& ptcl,
    thread uint& cmd_offset,
    thread uint& cmd_limit,
    constant _mslBufferSizes& _buffer_sizes
) {
    alloc_cmd(3u, config, bump, ptcl, cmd_offset, cmd_limit, _buffer_sizes);
    uint _e4 = cmd_offset;
    if (uint(_e4) < 1 + (_buffer_sizes.size8 - 0 - 4) / 4) {
        ptcl[_e4] = CMD_END_CLIP;
    }
    uint _e9 = cmd_offset;
    uint _e11 = _e9 + 1u;
    if (uint(_e11) < 1 + (_buffer_sizes.size8 - 0 - 4) / 4) {
        ptcl[_e11] = end_clip.blend;
    }
    uint _e16 = cmd_offset;
    uint _e18 = _e16 + 2u;
    if (uint(_e18) < 1 + (_buffer_sizes.size8 - 0 - 4) / 4) {
        ptcl[_e18] = as_type<uint>(end_clip.alpha);
    }
    uint _e24 = cmd_offset;
    cmd_offset = _e24 + 3u;
    return;
}

struct main_Input {
};
kernel void main_(
  metal::uint3 local_id [[thread_position_in_threadgroup]]
, metal::uint3 wg_id [[threadgroup_position_in_grid]]
, constant Config& config [[buffer(0)]]
, device type_7 const& scene [[buffer(1)]]
, device type_8 const& draw_monoids [[buffer(2)]]
, device type_9 const& bin_headers [[buffer(3)]]
, device type_7 const& info_bin_data [[buffer(4)]]
, device type_10 const& paths [[buffer(5)]]
, device type_11& tiles [[buffer(6)]]
, device BumpAllocators& bump [[buffer(7)]]
, device type_7& ptcl [[buffer(8)]]
, threadgroup type_13& sh_bitmaps
, threadgroup type_14& sh_part_count
, threadgroup type_14& sh_part_offsets
, threadgroup type_14& sh_drawobj_ix
, threadgroup type_14& sh_tile_stride
, threadgroup type_14& sh_tile_width
, threadgroup type_14& sh_tile_x0y0_
, threadgroup type_14& sh_tile_count
, threadgroup type_14& sh_tile_base
, constant _mslBufferSizes& _buffer_sizes [[buffer(9)]]
) {
    if (metal::all(local_id == metal::uint3(0u))) {
        for (int __i0 = 0; __i0 < 8; __i0++) {
            for (int __i1 = 0; __i1 < 256; __i1++) {
                metal::atomic_store_explicit(&sh_bitmaps.inner[__i0].inner[__i1], 0, metal::memory_order_relaxed);
            }
        }
        sh_part_count = {};
        sh_part_offsets = {};
        sh_drawobj_ix = {};
        sh_tile_stride = {};
        sh_tile_width = {};
        sh_tile_x0y0_ = {};
        sh_tile_count = {};
        sh_tile_base = {};
    }
    metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
    uint cmd_offset = {};
    uint cmd_limit = {};
    uint failed = {};
    uint clip_zero_depth = 0u;
    uint clip_depth = 0u;
    uint partition_ix = 0u;
    uint rd_ix = 0u;
    uint wr_ix = 0u;
    uint part_start_ix = 0u;
    uint ready_ix = 0u;
    uint render_blend_depth = 0u;
    uint max_blend_depth = 0u;
    uint i = {};
    uint count = {};
    uint i_1 = {};
    uint ix = {};
    uint part_ix = {};
    uint i_2 = {};
    uint tag = {};
    uint drawobj_ix = {};
    uint tile_count = {};
    uint i_3 = {};
    uint ix_1 = {};
    uint el_ix = {};
    uint i_4 = {};
    bool is_blend = {};
    uint slice_ix = {};
    uint bitmap = {};
    uint blend_ix = 0u;
    if (local_id.x == 0u) {
        uint _e7 = metal::atomic_load_explicit(&bump.failed, metal::memory_order_relaxed);
        failed = _e7 & 7u;
        uint _e13 = metal::atomic_load_explicit(&bump.seg_counts, metal::memory_order_relaxed);
        uint _e16 = config.seg_counts_size;
        if (_e13 > _e16) {
            uint _e19 = failed;
            failed = _e19 | STAGE_PATH_COUNT;
        }
        uint _e23 = failed;
        sh_part_count.inner[0] = static_cast<uint>(_e23);
    }
    metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
    uint unnamed = sh_part_count.inner[0];
    metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
    if (unnamed != 0u) {
        if ((wg_id.x == 0u) && (local_id.x == 0u)) {
            uint _e39 = metal::atomic_fetch_or_explicit(&bump.failed, unnamed, metal::memory_order_relaxed);
        }
        return;
    }
    uint _e42 = config.width_in_tiles;
    uint width_in_bins = ((_e42 + N_TILE_X) - 1u) / N_TILE_X;
    uint bin_ix = (width_in_bins * wg_id.y) + wg_id.x;
    uint _e55 = config.n_drawobj;
    uint n_partitions = ((_e55 + N_TILE) - 1u) / N_TILE;
    uint bin_tile_x = N_TILE_X * wg_id.x;
    uint bin_tile_y = N_TILE_Y * wg_id.y;
    uint tile_x = local_id.x % N_TILE_X;
    uint tile_y = local_id.x / N_TILE_X;
    uint _e77 = config.width_in_tiles;
    uint this_tile_ix = (((bin_tile_y + tile_y) * _e77) + bin_tile_x) + tile_x;
    cmd_offset = this_tile_ix * PTCL_INITIAL_ALLOC;
    uint _e86 = cmd_offset;
    cmd_limit = _e86 + 62u;
    uint blend_offset = cmd_offset;
    uint _e111 = cmd_offset;
    cmd_offset = _e111 + 1u;
    while(true) {
        if (true) {
        } else {
            break;
        }
        {
            i = 0u;
            bool loop_init = true;
            while(true) {
                if (!loop_init) {
                    uint _e126 = i;
                    i = _e126 + 1u;
                }
                loop_init = false;
                uint _e116 = i;
                if (_e116 < N_SLICE) {
                } else {
                    break;
                }
                {
                    uint _e120 = i;
                    uint _e122 = local_id.x;
                    if (uint(_e122) < 256 && uint(_e120) < 8) {
                        metal::atomic_store_explicit(&sh_bitmaps.inner[_e120].inner[_e122], 0u, metal::memory_order_relaxed);
                    }
                }
            }
            while(true) {
                if (true) {
                } else {
                    break;
                }
                {
                    uint _e129 = ready_ix;
                    uint _e130 = wr_ix;
                    uint _e132 = partition_ix;
                    if ((_e129 == _e130) && (_e132 < n_partitions)) {
                        uint _e135 = ready_ix;
                        part_start_ix = _e135;
                        count = 0u;
                        uint _e138 = partition_ix;
                        if ((_e138 + local_id.x) < n_partitions) {
                            uint _e142 = partition_ix;
                            uint in_ix = ((_e142 + local_id.x) * N_TILE) + bin_ix;
                            BinHeader bin_header = uint(in_ix) < 1 + (_buffer_sizes.size3 - 0 - 8) / 8 ? bin_headers[in_ix] : DefaultConstructible();
                            count = bin_header.element_count;
                            uint _e153 = local_id.x;
                            if (uint(_e153) < 256) {
                                sh_part_offsets.inner[_e153] = bin_header.chunk_offset;
                            }
                        }
                        i_1 = 0u;
                        bool loop_init_1 = true;
                        while(true) {
                            if (!loop_init_1) {
                                uint _e182 = i_1;
                                i_1 = _e182 + 1u;
                            }
                            loop_init_1 = false;
                            uint _e158 = i_1;
                            if (_e158 < (((metal::ctz(WG_SIZE) + 1) % 33) - 1)) {
                            } else {
                                break;
                            }
                            {
                                uint _e163 = local_id.x;
                                uint _e165 = count;
                                if (uint(_e163) < 256) {
                                    sh_part_count.inner[_e163] = _e165;
                                }
                                metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
                                uint _e168 = i_1;
                                if (local_id.x >= (1u << _e168)) {
                                    uint _e174 = i_1;
                                    uint _e176 = local_id.x - (1u << _e174);
                                    uint _e178 = uint(_e176) < 256 ? sh_part_count.inner[_e176] : DefaultConstructible();
                                    uint _e179 = count;
                                    count = _e179 + _e178;
                                }
                                metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
                            }
                        }
                        uint _e185 = local_id.x;
                        uint _e187 = part_start_ix;
                        uint _e188 = count;
                        if (uint(_e185) < 256) {
                            sh_part_count.inner[_e185] = _e187 + _e188;
                        }
                        metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
                        uint unnamed_1 = sh_part_count.inner[255];
                        metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
                        ready_ix = unnamed_1;
                        uint _e194 = partition_ix;
                        partition_ix = _e194 + WG_SIZE;
                    }
                    uint _e196 = rd_ix;
                    ix = _e196 + local_id.x;
                    uint _e200 = ix;
                    uint _e201 = wr_ix;
                    uint _e203 = ix;
                    uint _e204 = ready_ix;
                    if ((_e200 >= _e201) && (_e203 < _e204)) {
                        part_ix = 0u;
                        i_2 = 0u;
                        bool loop_init_2 = true;
                        while(true) {
                            if (!loop_init_2) {
                                uint _e228 = i_2;
                                i_2 = _e228 + 1u;
                            }
                            loop_init_2 = false;
                            uint _e211 = i_2;
                            if (_e211 < (((metal::ctz(WG_SIZE) + 1) % 33) - 1)) {
                            } else {
                                break;
                            }
                            {
                                uint _e215 = part_ix;
                                uint _e217 = i_2;
                                uint probe = _e215 + (128u >> _e217);
                                uint _e220 = ix;
                                uint _e223 = probe - 1u;
                                uint _e225 = uint(_e223) < 256 ? sh_part_count.inner[_e223] : DefaultConstructible();
                                if (_e220 >= _e225) {
                                    part_ix = probe;
                                }
                            }
                        }
                        uint _e230 = part_start_ix;
                        uint _e232 = part_ix;
                        uint _e234 = _e232 - 1u;
                        uint _e236 = uint(_e234) < 256 ? sh_part_count.inner[_e234] : DefaultConstructible();
                        uint _e237 = part_ix;
                        uint _e241 = ix;
                        ix = _e241 - ((_e237 > 0u) ? _e236 : _e230);
                        uint _e245 = config.bin_data_start;
                        uint _e247 = part_ix;
                        uint _e249 = uint(_e247) < 256 ? sh_part_offsets.inner[_e247] : DefaultConstructible();
                        uint offset = _e245 + _e249;
                        uint _e252 = local_id.x;
                        uint _e255 = ix;
                        uint _e256 = offset + _e255;
                        uint _e258 = uint(_e256) < 1 + (_buffer_sizes.size4 - 0 - 4) / 4 ? info_bin_data[_e256] : DefaultConstructible();
                        if (uint(_e252) < 256) {
                            sh_drawobj_ix.inner[_e252] = _e258;
                        }
                    }
                    uint _e259 = rd_ix;
                    uint _e262 = ready_ix;
                    wr_ix = metal::min(_e259 + N_TILE, _e262);
                    uint _e264 = wr_ix;
                    uint _e265 = rd_ix;
                    uint _e269 = wr_ix;
                    uint _e270 = ready_ix;
                    uint _e272 = partition_ix;
                    if (((_e264 - _e265) >= N_TILE) || ((_e269 >= _e270) && (_e272 >= n_partitions))) {
                        break;
                    }
                    metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
                }
            }
            tag = DRAWTAG_NOP;
            uint _e280 = rd_ix;
            uint _e282 = wr_ix;
            if ((local_id.x + _e280) < _e282) {
                uint _e285 = local_id.x;
                uint _e287 = uint(_e285) < 256 ? sh_drawobj_ix.inner[_e285] : DefaultConstructible();
                drawobj_ix = _e287;
                uint _e291 = config.drawtag_base;
                uint _e292 = drawobj_ix;
                uint _e293 = _e291 + _e292;
                uint _e295 = uint(_e293) < 1 + (_buffer_sizes.size1 - 0 - 4) / 4 ? scene[_e293] : DefaultConstructible();
                tag = _e295;
            }
            tile_count = 0u;
            uint _e298 = tag;
            if (_e298 != DRAWTAG_NOP) {
                uint _e302 = drawobj_ix;
                uint path_ix = uint(_e302) < 1 + (_buffer_sizes.size2 - 0 - 16) / 16 ? draw_monoids[_e302].path_ix : DefaultConstructible();
                Path path = uint(path_ix) < 1 + (_buffer_sizes.size5 - 0 - 32) / 32 ? paths[path_ix] : DefaultConstructible();
                uint stride = path.bbox.z - path.bbox.x;
                uint _e315 = local_id.x;
                if (uint(_e315) < 256) {
                    sh_tile_stride.inner[_e315] = stride;
                }
                int dx = static_cast<int>(path.bbox.x) - static_cast<int>(bin_tile_x);
                int dy = static_cast<int>(path.bbox.y) - static_cast<int>(bin_tile_y);
                int x0_ = metal::clamp(dx, 0, 16);
                int y0_ = metal::clamp(dy, 0, 16);
                int x1_ = metal::clamp(static_cast<int>(path.bbox.z) - static_cast<int>(bin_tile_x), 0, 16);
                int y1_ = metal::clamp(static_cast<int>(path.bbox.w) - static_cast<int>(bin_tile_y), 0, 16);
                uint _e350 = local_id.x;
                if (uint(_e350) < 256) {
                    sh_tile_width.inner[_e350] = static_cast<uint>(x1_ - x0_);
                }
                uint _e355 = local_id.x;
                if (uint(_e355) < 256) {
                    sh_tile_x0y0_.inner[_e355] = static_cast<uint>(x0_) | static_cast<uint>(y0_ << 16u);
                }
                tile_count = static_cast<uint>(x1_ - x0_) * static_cast<uint>(y1_ - y0_);
                uint base = path.tiles - static_cast<uint>((dy * static_cast<int>(stride)) + dx);
                uint _e374 = local_id.x;
                if (uint(_e374) < 256) {
                    sh_tile_base.inner[_e374] = base;
                }
            }
            uint _e377 = local_id.x;
            uint _e379 = tile_count;
            if (uint(_e377) < 256) {
                sh_tile_count.inner[_e377] = _e379;
            }
            i_3 = 0u;
            bool loop_init_3 = true;
            while(true) {
                if (!loop_init_3) {
                    uint _e406 = i_3;
                    i_3 = _e406 + 1u;
                }
                loop_init_3 = false;
                uint _e382 = i_3;
                if (_e382 < (((metal::ctz(N_TILE) + 1) % 33) - 1)) {
                } else {
                    break;
                }
                {
                    metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
                    uint _e388 = i_3;
                    if (local_id.x >= (1u << _e388)) {
                        uint _e394 = i_3;
                        uint _e396 = local_id.x - (1u << _e394);
                        uint _e398 = uint(_e396) < 256 ? sh_tile_count.inner[_e396] : DefaultConstructible();
                        uint _e399 = tile_count;
                        tile_count = _e399 + _e398;
                    }
                    metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
                    uint _e402 = local_id.x;
                    uint _e404 = tile_count;
                    if (uint(_e402) < 256) {
                        sh_tile_count.inner[_e402] = _e404;
                    }
                }
            }
            metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
            uint total_tile_count = sh_tile_count.inner[255];
            ix_1 = local_id.x;
            bool loop_init_4 = true;
            while(true) {
                if (!loop_init_4) {
                    uint _e557 = ix_1;
                    ix_1 = _e557 + N_TILE;
                }
                loop_init_4 = false;
                uint _e413 = ix_1;
                if (_e413 < total_tile_count) {
                } else {
                    break;
                }
                {
                    el_ix = 0u;
                    i_4 = 0u;
                    bool loop_init_5 = true;
                    while(true) {
                        if (!loop_init_5) {
                            uint _e436 = i_4;
                            i_4 = _e436 + 1u;
                        }
                        loop_init_5 = false;
                        uint _e419 = i_4;
                        if (_e419 < (((metal::ctz(N_TILE) + 1) % 33) - 1)) {
                        } else {
                            break;
                        }
                        {
                            uint _e423 = el_ix;
                            uint _e425 = i_4;
                            uint probe_1 = _e423 + (128u >> _e425);
                            uint _e428 = ix_1;
                            uint _e431 = probe_1 - 1u;
                            uint _e433 = uint(_e431) < 256 ? sh_tile_count.inner[_e431] : DefaultConstructible();
                            if (_e428 >= _e433) {
                                el_ix = probe_1;
                            }
                        }
                    }
                    uint _e439 = el_ix;
                    uint _e441 = uint(_e439) < 256 ? sh_drawobj_ix.inner[_e439] : DefaultConstructible();
                    drawobj_ix = _e441;
                    uint _e445 = config.drawtag_base;
                    uint _e446 = drawobj_ix;
                    uint _e447 = _e445 + _e446;
                    uint _e449 = uint(_e447) < 1 + (_buffer_sizes.size1 - 0 - 4) / 4 ? scene[_e447] : DefaultConstructible();
                    tag = _e449;
                    uint _e450 = ix_1;
                    uint _e453 = el_ix;
                    uint _e455 = _e453 - 1u;
                    uint _e457 = uint(_e455) < 256 ? sh_tile_count.inner[_e455] : DefaultConstructible();
                    uint _e458 = el_ix;
                    uint seq_ix = _e450 - ((_e458 > 0u) ? _e457 : 0u);
                    uint _e464 = el_ix;
                    uint width = uint(_e464) < 256 ? sh_tile_width.inner[_e464] : DefaultConstructible();
                    uint _e468 = el_ix;
                    uint x0y0_ = uint(_e468) < 256 ? sh_tile_x0y0_.inner[_e468] : DefaultConstructible();
                    uint x = (x0y0_ & 65535u) + (seq_ix % width);
                    uint y = (x0y0_ >> 16u) + (seq_ix / width);
                    uint _e480 = el_ix;
                    uint _e482 = uint(_e480) < 256 ? sh_tile_base.inner[_e480] : DefaultConstructible();
                    uint _e484 = el_ix;
                    uint _e486 = uint(_e484) < 256 ? sh_tile_stride.inner[_e484] : DefaultConstructible();
                    uint tile_ix_1 = (_e482 + (_e486 * y)) + x;
                    Tile tile_1 = uint(tile_ix_1) < 1 + (_buffer_sizes.size6 - 0 - 8) / 8 ? tiles[tile_ix_1] : DefaultConstructible();
                    uint _e493 = tag;
                    bool is_clip = (_e493 & 1u) != 0u;
                    is_blend = false;
                    if (is_clip) {
                        uint _e502 = drawobj_ix;
                        uint scene_offset = uint(_e502) < 1 + (_buffer_sizes.size2 - 0 - 16) / 16 ? draw_monoids[_e502].scene_offset : DefaultConstructible();
                        uint _e508 = config.drawdata_base;
                        uint dd = _e508 + scene_offset;
                        uint blend = uint(dd) < 1 + (_buffer_sizes.size1 - 0 - 4) / 4 ? scene[dd] : DefaultConstructible();
                        is_blend = blend != 32771u;
                    }
                    uint _e515 = drawobj_ix;
                    uint di = uint(_e515) < 1 + (_buffer_sizes.size2 - 0 - 16) / 16 ? draw_monoids[_e515].info_offset : DefaultConstructible();
                    uint draw_flags_1 = uint(di) < 1 + (_buffer_sizes.size4 - 0 - 4) / 4 ? info_bin_data[di] : DefaultConstructible();
                    bool even_odd_1 = (draw_flags_1 & DRAW_INFO_FLAGS_FILL_RULE_BIT) != 0u;
                    uint n_segs_1 = tile_1.segment_count_or_ix;
                    bool backdrop_clear = (even_odd_1 ? (metal::abs(tile_1.backdrop) & 1) : tile_1.backdrop) == 0;
                    bool _e539 = is_blend;
                    bool include_tile = ((n_segs_1 != 0u) || (backdrop_clear == is_clip)) || _e539;
                    if (include_tile) {
                        uint _e541 = el_ix;
                        uint el_slice = _e541 / 32u;
                        uint _e545 = el_ix;
                        uint el_mask = 1u << (_e545 & 31u);
                        uint _e553 = (y * N_TILE_X) + x;
                        uint _e555 = uint(_e553) < 256 && uint(el_slice) < 8 ? metal::atomic_fetch_or_explicit(&sh_bitmaps.inner[el_slice].inner[_e553], el_mask, metal::memory_order_relaxed) : DefaultConstructible();
                    }
                }
            }
            metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
            slice_ix = 0u;
            uint _e563 = local_id.x;
            uint _e565 = uint(_e563) < 256 ? metal::atomic_load_explicit(&sh_bitmaps.inner[0].inner[_e563], metal::memory_order_relaxed) : DefaultConstructible();
            bitmap = _e565;
            while(true) {
                if (true) {
                } else {
                    break;
                }
                {
                    uint _e568 = bitmap;
                    if (_e568 == 0u) {
                        uint _e572 = slice_ix;
                        slice_ix = _e572 + 1u;
                        uint _e574 = slice_ix;
                        if (_e574 == N_SLICE) {
                            break;
                        }
                        uint _e578 = slice_ix;
                        uint _e580 = local_id.x;
                        uint _e582 = uint(_e580) < 256 && uint(_e578) < 8 ? metal::atomic_load_explicit(&sh_bitmaps.inner[_e578].inner[_e580], metal::memory_order_relaxed) : DefaultConstructible();
                        bitmap = _e582;
                        uint _e583 = bitmap;
                        if (_e583 == 0u) {
                            continue;
                        }
                    }
                    uint _e586 = slice_ix;
                    uint _e589 = bitmap;
                    uint el_ix_1 = (_e586 * 32u) + (((metal::ctz(_e589) + 1) % 33) - 1);
                    uint _e594 = uint(el_ix_1) < 256 ? sh_drawobj_ix.inner[el_ix_1] : DefaultConstructible();
                    drawobj_ix = _e594;
                    uint _e595 = bitmap;
                    uint _e598 = bitmap;
                    bitmap = _e598 & (_e595 - 1u);
                    uint _e603 = config.drawtag_base;
                    uint _e604 = drawobj_ix;
                    uint _e605 = _e603 + _e604;
                    uint drawtag = uint(_e605) < 1 + (_buffer_sizes.size1 - 0 - 4) / 4 ? scene[_e605] : DefaultConstructible();
                    uint _e609 = drawobj_ix;
                    DrawMonoid dm = uint(_e609) < 1 + (_buffer_sizes.size2 - 0 - 16) / 16 ? draw_monoids[_e609] : DefaultConstructible();
                    uint _e614 = config.drawdata_base;
                    uint dd_1 = _e614 + dm.scene_offset;
                    uint di_1 = dm.info_offset;
                    uint draw_flags_2 = uint(di_1) < 1 + (_buffer_sizes.size4 - 0 - 4) / 4 ? info_bin_data[di_1] : DefaultConstructible();
                    uint _e621 = clip_zero_depth;
                    if (_e621 == 0u) {
                        uint _e626 = uint(el_ix_1) < 256 ? sh_tile_base.inner[el_ix_1] : DefaultConstructible();
                        uint _e629 = uint(el_ix_1) < 256 ? sh_tile_stride.inner[el_ix_1] : DefaultConstructible();
                        uint tile_ix_2 = (_e626 + (_e629 * tile_y)) + tile_x;
                        Tile tile_2 = uint(tile_ix_2) < 1 + (_buffer_sizes.size6 - 0 - 8) / 8 ? tiles[tile_ix_2] : DefaultConstructible();
                        switch(drawtag) {
                            case 68u: {
                                write_path(tile_2, tile_ix_2, draw_flags_2, config, tiles, bump, ptcl, cmd_offset, cmd_limit, _buffer_sizes);
                                uint rgba_color = uint(dd_1) < 1 + (_buffer_sizes.size1 - 0 - 4) / 4 ? scene[dd_1] : DefaultConstructible();
                                write_color(CmdColor {rgba_color}, config, bump, ptcl, cmd_offset, cmd_limit, _buffer_sizes);
                                break;
                            }
                            case 276u: {
                                write_path(tile_2, tile_ix_2, draw_flags_2, config, tiles, bump, ptcl, cmd_offset, cmd_limit, _buffer_sizes);
                                uint index_1 = uint(dd_1) < 1 + (_buffer_sizes.size1 - 0 - 4) / 4 ? scene[dd_1] : DefaultConstructible();
                                uint info_offset_2 = di_1 + 1u;
                                write_grad(CMD_LIN_GRAD, index_1, info_offset_2, config, bump, ptcl, cmd_offset, cmd_limit, _buffer_sizes);
                                break;
                            }
                            case 668u: {
                                write_path(tile_2, tile_ix_2, draw_flags_2, config, tiles, bump, ptcl, cmd_offset, cmd_limit, _buffer_sizes);
                                uint index_2 = uint(dd_1) < 1 + (_buffer_sizes.size1 - 0 - 4) / 4 ? scene[dd_1] : DefaultConstructible();
                                uint info_offset_3 = di_1 + 1u;
                                write_grad(CMD_RAD_GRAD, index_2, info_offset_3, config, bump, ptcl, cmd_offset, cmd_limit, _buffer_sizes);
                                break;
                            }
                            case 596u: {
                                write_path(tile_2, tile_ix_2, draw_flags_2, config, tiles, bump, ptcl, cmd_offset, cmd_limit, _buffer_sizes);
                                uint index_3 = uint(dd_1) < 1 + (_buffer_sizes.size1 - 0 - 4) / 4 ? scene[dd_1] : DefaultConstructible();
                                uint info_offset_4 = di_1 + 1u;
                                write_grad(CMD_SWEEP_GRAD, index_3, info_offset_4, config, bump, ptcl, cmd_offset, cmd_limit, _buffer_sizes);
                                break;
                            }
                            case 584u: {
                                write_path(tile_2, tile_ix_2, draw_flags_2, config, tiles, bump, ptcl, cmd_offset, cmd_limit, _buffer_sizes);
                                write_image(di_1 + 1u, config, bump, ptcl, cmd_offset, cmd_limit, _buffer_sizes);
                                break;
                            }
                            case 9u: {
                                if ((tile_2.segment_count_or_ix == 0u) && (tile_2.backdrop == 0)) {
                                    uint _e667 = clip_depth;
                                    clip_zero_depth = _e667 + 1u;
                                } else {
                                    write_begin_clip(config, bump, ptcl, cmd_offset, cmd_limit, _buffer_sizes);
                                    uint _e671 = render_blend_depth;
                                    render_blend_depth = _e671 + 1u;
                                    uint _e673 = max_blend_depth;
                                    uint _e674 = render_blend_depth;
                                    max_blend_depth = metal::max(_e673, _e674);
                                }
                                uint _e677 = clip_depth;
                                clip_depth = _e677 + 1u;
                                break;
                            }
                            case 33u: {
                                uint _e680 = clip_depth;
                                clip_depth = _e680 - 1u;
                                write_path(tile_2, tile_ix_2, 0u, config, tiles, bump, ptcl, cmd_offset, cmd_limit, _buffer_sizes);
                                uint blend_1 = uint(dd_1) < 1 + (_buffer_sizes.size1 - 0 - 4) / 4 ? scene[dd_1] : DefaultConstructible();
                                uint _e688 = dd_1 + 1u;
                                uint _e690 = uint(_e688) < 1 + (_buffer_sizes.size1 - 0 - 4) / 4 ? scene[_e688] : DefaultConstructible();
                                float alpha = as_type<float>(_e690);
                                write_end_clip(CmdEndClip {blend_1, alpha}, config, bump, ptcl, cmd_offset, cmd_limit, _buffer_sizes);
                                uint _e694 = render_blend_depth;
                                render_blend_depth = _e694 - 1u;
                                break;
                            }
                            default: {
                                break;
                            }
                        }
                    } else {
                        switch(drawtag) {
                            case 9u: {
                                uint _e697 = clip_depth;
                                clip_depth = _e697 + 1u;
                                break;
                            }
                            case 33u: {
                                uint _e699 = clip_depth;
                                uint _e700 = clip_zero_depth;
                                if (_e699 == _e700) {
                                    clip_zero_depth = 0u;
                                }
                                uint _e704 = clip_depth;
                                clip_depth = _e704 - 1u;
                                break;
                            }
                            default: {
                                break;
                            }
                        }
                    }
                }
            }
            uint _e707 = rd_ix;
            rd_ix = _e707 + N_TILE;
            uint _e709 = rd_ix;
            uint _e710 = ready_ix;
            uint _e712 = partition_ix;
            if ((_e709 >= _e710) && (_e712 >= n_partitions)) {
                break;
            }
            metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
        }
    }
    uint _e718 = config.width_in_tiles;
    uint _e723 = config.height_in_tiles;
    if (((bin_tile_x + tile_x) < _e718) && ((bin_tile_y + tile_y) < _e723)) {
        uint _e728 = cmd_offset;
        if (uint(_e728) < 1 + (_buffer_sizes.size8 - 0 - 4) / 4) {
            ptcl[_e728] = CMD_END;
        }
        uint _e733 = max_blend_depth;
        if (_e733 > BLEND_STACK_SPLIT) {
            uint _e736 = max_blend_depth;
            uint scratch_size = ((_e736 - BLEND_STACK_SPLIT) * TILE_WIDTH) * TILE_HEIGHT;
            uint _e745 = metal::atomic_fetch_add_explicit(&bump.blend, scratch_size, metal::memory_order_relaxed);
            blend_ix = _e745;
            uint _e746 = blend_ix;
            uint _e750 = config.blend_size;
            if ((_e746 + scratch_size) > _e750) {
                uint _e755 = metal::atomic_fetch_or_explicit(&bump.failed, STAGE_COARSE, metal::memory_order_relaxed);
            }
        }
        uint _e758 = blend_ix;
        if (uint(blend_offset) < 1 + (_buffer_sizes.size8 - 0 - 4) / 4) {
            ptcl[blend_offset] = _e758;
        }
        return;
    } else {
        return;
    }
}