// language: metal2.0 #include #include using metal::uint; 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(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; ptcl[_e36] = CMD_JUMP; uint _e41 = cmd_offset; uint _e45 = new_cmd; ptcl[_e41 + 1u] = _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; tiles[tile_ix].segment_count_or_ix = ~(_e13); alloc_cmd(4u, config, bump, ptcl, cmd_offset, cmd_limit, _buffer_sizes); uint _e18 = cmd_offset; 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(even_odd); uint _e29 = seg_ix; CmdFill fill = CmdFill {size_and_rule, _e29, tile.backdrop}; uint _e34 = cmd_offset; ptcl[_e34 + 1u] = fill.size_and_rule; uint _e41 = cmd_offset; ptcl[_e41 + 2u] = fill.seg_data; uint _e48 = cmd_offset; ptcl[_e48 + 3u] = static_cast(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; 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; ptcl[_e4] = CMD_COLOR; uint _e9 = cmd_offset; ptcl[_e9 + 1u] = 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; ptcl[_e6] = ty; uint _e10 = cmd_offset; ptcl[_e10 + 1u] = index; uint _e16 = cmd_offset; ptcl[_e16 + 2u] = 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; ptcl[_e4] = CMD_IMAGE; uint _e9 = cmd_offset; ptcl[_e9 + 1u] = 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; 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; ptcl[_e4] = CMD_END_CLIP; uint _e9 = cmd_offset; ptcl[_e9 + 1u] = end_clip.blend; uint _e16 = cmd_offset; ptcl[_e16 + 2u] = as_type(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(30)]] ) { 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(_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; metal::atomic_store_explicit(&sh_bitmaps.inner[_e120].inner[local_id.x], 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 = bin_headers[in_ix]; count = bin_header.element_count; sh_part_offsets.inner[local_id.x] = 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 _e165 = count; sh_part_count.inner[local_id.x] = _e165; metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); uint _e168 = i_1; if (local_id.x >= (1u << _e168)) { uint _e174 = i_1; uint _e178 = sh_part_count.inner[local_id.x - (1u << _e174)]; uint _e179 = count; count = _e179 + _e178; } metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); } } uint _e187 = part_start_ix; uint _e188 = count; sh_part_count.inner[local_id.x] = _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 _e225 = sh_part_count.inner[probe - 1u]; if (_e220 >= _e225) { part_ix = probe; } } } uint _e230 = part_start_ix; uint _e232 = part_ix; uint _e236 = sh_part_count.inner[_e232 - 1u]; 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 = sh_part_offsets.inner[_e247]; uint offset = _e245 + _e249; uint _e255 = ix; uint _e258 = info_bin_data[offset + _e255]; sh_drawobj_ix.inner[local_id.x] = _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 _e287 = sh_drawobj_ix.inner[local_id.x]; drawobj_ix = _e287; uint _e291 = config.drawtag_base; uint _e292 = drawobj_ix; uint _e295 = scene[_e291 + _e292]; tag = _e295; } tile_count = 0u; uint _e298 = tag; if (_e298 != DRAWTAG_NOP) { uint _e302 = drawobj_ix; uint path_ix = draw_monoids[_e302].path_ix; Path path = paths[path_ix]; uint stride = path.bbox.z - path.bbox.x; sh_tile_stride.inner[local_id.x] = stride; int dx = static_cast(path.bbox.x) - static_cast(bin_tile_x); int dy = static_cast(path.bbox.y) - static_cast(bin_tile_y); int x0_ = metal::clamp(dx, 0, 16); int y0_ = metal::clamp(dy, 0, 16); int x1_ = metal::clamp(static_cast(path.bbox.z) - static_cast(bin_tile_x), 0, 16); int y1_ = metal::clamp(static_cast(path.bbox.w) - static_cast(bin_tile_y), 0, 16); sh_tile_width.inner[local_id.x] = static_cast(x1_ - x0_); sh_tile_x0y0_.inner[local_id.x] = static_cast(x0_) | static_cast(y0_ << 16u); tile_count = static_cast(x1_ - x0_) * static_cast(y1_ - y0_); uint base = path.tiles - static_cast((dy * static_cast(stride)) + dx); sh_tile_base.inner[local_id.x] = base; } uint _e379 = tile_count; sh_tile_count.inner[local_id.x] = _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 _e398 = sh_tile_count.inner[local_id.x - (1u << _e394)]; uint _e399 = tile_count; tile_count = _e399 + _e398; } metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); uint _e404 = tile_count; sh_tile_count.inner[local_id.x] = _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 _e433 = sh_tile_count.inner[probe_1 - 1u]; if (_e428 >= _e433) { el_ix = probe_1; } } } uint _e439 = el_ix; uint _e441 = sh_drawobj_ix.inner[_e439]; drawobj_ix = _e441; uint _e445 = config.drawtag_base; uint _e446 = drawobj_ix; uint _e449 = scene[_e445 + _e446]; tag = _e449; uint _e450 = ix_1; uint _e453 = el_ix; uint _e457 = sh_tile_count.inner[_e453 - 1u]; uint _e458 = el_ix; uint seq_ix = _e450 - ((_e458 > 0u) ? _e457 : 0u); uint _e464 = el_ix; uint width = sh_tile_width.inner[_e464]; uint _e468 = el_ix; uint x0y0_ = sh_tile_x0y0_.inner[_e468]; uint x = (x0y0_ & 65535u) + (seq_ix % width); uint y = (x0y0_ >> 16u) + (seq_ix / width); uint _e480 = el_ix; uint _e482 = sh_tile_base.inner[_e480]; uint _e484 = el_ix; uint _e486 = sh_tile_stride.inner[_e484]; uint tile_ix_1 = (_e482 + (_e486 * y)) + x; Tile tile_1 = tiles[tile_ix_1]; uint _e493 = tag; bool is_clip = (_e493 & 1u) != 0u; is_blend = false; if (is_clip) { uint _e502 = drawobj_ix; uint scene_offset = draw_monoids[_e502].scene_offset; uint _e508 = config.drawdata_base; uint dd = _e508 + scene_offset; uint blend = scene[dd]; is_blend = blend != 32771u; } uint _e515 = drawobj_ix; uint di = draw_monoids[_e515].info_offset; uint draw_flags_1 = info_bin_data[di]; 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 _e555 = metal::atomic_fetch_or_explicit(&sh_bitmaps.inner[el_slice].inner[(y * N_TILE_X) + x], el_mask, metal::memory_order_relaxed); } } } metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); slice_ix = 0u; uint _e565 = metal::atomic_load_explicit(&sh_bitmaps.inner[0].inner[local_id.x], metal::memory_order_relaxed); 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 _e582 = metal::atomic_load_explicit(&sh_bitmaps.inner[_e578].inner[local_id.x], metal::memory_order_relaxed); 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 = sh_drawobj_ix.inner[el_ix_1]; drawobj_ix = _e594; uint _e595 = bitmap; uint _e598 = bitmap; bitmap = _e598 & (_e595 - 1u); uint _e603 = config.drawtag_base; uint _e604 = drawobj_ix; uint drawtag = scene[_e603 + _e604]; uint _e609 = drawobj_ix; DrawMonoid dm = draw_monoids[_e609]; uint _e614 = config.drawdata_base; uint dd_1 = _e614 + dm.scene_offset; uint di_1 = dm.info_offset; uint draw_flags_2 = info_bin_data[di_1]; uint _e621 = clip_zero_depth; if (_e621 == 0u) { uint _e626 = sh_tile_base.inner[el_ix_1]; uint _e629 = sh_tile_stride.inner[el_ix_1]; uint tile_ix_2 = (_e626 + (_e629 * tile_y)) + tile_x; Tile tile_2 = tiles[tile_ix_2]; 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 = scene[dd_1]; 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 = scene[dd_1]; 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 = scene[dd_1]; 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 = scene[dd_1]; 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 = scene[dd_1]; uint _e690 = scene[dd_1 + 1u]; float alpha = as_type(_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; 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; ptcl[blend_offset] = _e758; return; } else { return; } }