Skip to content

Instantly share code, notes, and snippets.

@raphlinus
Created October 7, 2022 01:41
Show Gist options
  • Save raphlinus/1406c2a41589ea1a616cfc2552d6dc1c to your computer and use it in GitHub Desktop.
Save raphlinus/1406c2a41589ea1a616cfc2552d6dc1c to your computer and use it in GitHub Desktop.
Partially translated coarse.comp to wgsl
struct Alloc {
offset : u32,
}
struct BinInstanceRef {
offset : u32,
}
struct BinInstance {
element_ix : u32,
}
struct PathRef {
offset : u32,
}
struct TileRef {
offset : u32,
}
struct CmdRef {
offset : u32,
}
struct Path {
bbox : vec4<u32>,
tiles : TileRef,
}
struct TileSegRef {
offset : u32,
}
struct Tile {
tile : TileSegRef,
backdrop : i32,
}
struct CmdStrokeRef {
offset : u32,
}
struct CmdStroke {
tile_ref : u32,
half_width : f32,
}
struct CmdFillRef {
offset : u32,
}
struct CmdFill {
tile_ref : u32,
backdrop : i32,
}
struct CmdColorRef {
offset : u32,
}
struct CmdColor {
rgba_color : u32,
}
struct CmdLinGradRef {
offset : u32,
}
struct CmdLinGrad {
index : u32,
line_x : f32,
line_y : f32,
line_c : f32,
}
struct CmdRadGradRef {
offset : u32,
}
struct CmdRadGrad {
index : u32,
mat_ : vec4<f32>,
xlat : vec2<f32>,
c1 : vec2<f32>,
ra : f32,
roff : f32,
}
struct CmdImageRef {
offset : u32,
}
struct CmdImage {
index : u32,
offset : vec2<i32>,
}
struct CmdEndClipRef {
offset : u32,
}
struct CmdEndClip {
blend : u32,
}
struct CmdJumpRef {
offset : u32,
}
struct CmdJump {
new_ref : u32,
}
type RTArr = array<u32>;
type RTArr_1 = array<u32>;
struct Memory {
mem_offset : u32,
mem_error : u32,
blend_offset : u32,
memory : RTArr_1,
}
struct Alloc_1 {
offset : u32,
}
struct Config {
mem_size : u32,
n_elements : u32,
n_pathseg : u32,
width_in_tiles : u32,
height_in_tiles : u32,
tile_alloc : Alloc_1,
bin_alloc : Alloc_1,
ptcl_alloc : Alloc_1,
pathseg_alloc : Alloc_1,
anno_alloc : Alloc_1,
path_bbox_alloc : Alloc_1,
drawmonoid_alloc : Alloc_1,
clip_alloc : Alloc_1,
clip_bic_alloc : Alloc_1,
clip_stack_alloc : Alloc_1,
clip_bbox_alloc : Alloc_1,
draw_bbox_alloc : Alloc_1,
drawinfo_alloc : Alloc_1,
n_trans : u32,
n_path : u32,
n_clip : u32,
trans_offset : u32,
linewidth_offset : u32,
pathtag_offset : u32,
pathseg_offset : u32,
drawtag_offset : u32,
drawdata_offset : u32,
}
struct ConfigBuf {
conf : Config,
}
struct SceneBuf {
scene : RTArr_1,
}
var<private> mem_ok_2 : bool;
@group(0) @binding(0) var<storage, read_write> x_270 : Memory;
@group(0) @binding(1) var<storage, read> x_883 : ConfigBuf;
var<private> gl_WorkGroupID : vec3<u32>;
var<private> gl_LocalInvocationID : vec3<u32>;
var<workgroup> sh_bitmaps : array<array<u32, 256u>, 8u>;
var<workgroup> sh_part_elements : array<Alloc, 256u>;
var<workgroup> sh_part_count : array<u32, 256u>;
var<workgroup> sh_elements : array<u32, 256u>;
@group(0) @binding(2) var<storage, read> x_1382 : SceneBuf;
var<workgroup> sh_tile_stride : array<u32, 256u>;
var<workgroup> sh_tile_width : array<u32, 256u>;
var<workgroup> sh_tile_x0 : array<u32, 256u>;
var<workgroup> sh_tile_y0 : array<u32, 256u>;
var<workgroup> sh_tile_base : array<u32, 256u>;
var<workgroup> sh_tile_count : array<u32, 256u>;
fn check_deps_u1_(dep_stage : ptr<function, u32>) -> bool {
return true;
}
fn slice_mem_struct_Alloc_u11_u1_u1_(a : ptr<function, Alloc>, offset_5 : ptr<function, u32>, size_2 : ptr<function, u32>) -> Alloc {
let x_319 : u32 = (*(a)).offset;
let x_320 : u32 = *(offset_5);
return Alloc((x_319 + x_320));
}
fn touch_mem_struct_Alloc_u11_u1_(alloc : ptr<function, Alloc>, offset_2 : ptr<function, u32>) -> bool {
return true;
}
fn read_mem_struct_Alloc_u11_u1_(alloc_2 : ptr<function, Alloc>, offset_4 : ptr<function, u32>) -> u32 {
var param_2 : Alloc;
var param_3 : u32;
var v : u32;
let x_303 : Alloc = *(alloc_2);
param_2 = x_303;
let x_305 : u32 = *(offset_4);
param_3 = x_305;
let x_306 : bool = touch_mem_struct_Alloc_u11_u1_(&(param_2), &(param_3));
if (!(x_306)) {
return 0u;
}
let x_312 : u32 = *(offset_4);
let x_314 : u32 = x_270.memory[x_312];
v = x_314;
let x_315 : u32 = v;
return x_315;
}
fn new_alloc_u1_u1_b1_(offset_1 : ptr<function, u32>, size : ptr<function, u32>, mem_ok : ptr<function, bool>) -> Alloc {
var a_24 : Alloc;
let x_259 : u32 = *(offset_1);
a_24.offset = x_259;
let x_261 : Alloc = a_24;
return x_261;
}
fn BinInstance_index_struct_BinInstanceRef_u11_u1_(reff : ptr<function, BinInstanceRef>, index : ptr<function, u32>) -> BinInstanceRef {
let x_326 : u32 = (*(reff)).offset;
let x_327 : u32 = *(index);
return BinInstanceRef((x_326 + (x_327 * 4u)));
}
fn BinInstance_read_struct_Alloc_u11_struct_BinInstanceRef_u11_(a_1 : ptr<function, Alloc>, ref_1 : ptr<function, BinInstanceRef>) -> BinInstance {
var ix : u32;
var raw0 : u32;
var param_4 : Alloc;
var param_5 : u32;
var s_16 : BinInstance;
let x_336 : u32 = (*(ref_1)).offset;
ix = (x_336 >> bitcast<u32>(2i));
let x_340 : u32 = ix;
let x_343 : Alloc = *(a_1);
param_4 = x_343;
param_5 = (x_340 + 0u);
let x_345 : u32 = read_mem_struct_Alloc_u11_u1_(&(param_4), &(param_5));
raw0 = x_345;
let x_348 : u32 = raw0;
s_16.element_ix = x_348;
let x_350 : BinInstance = s_16;
return x_350;
}
fn Path_read_struct_Alloc_u11_struct_PathRef_u11_(a_2 : ptr<function, Alloc>, ref_2 : ptr<function, PathRef>) -> Path {
var ix_1 : u32;
var raw0_1 : u32;
var param_6 : Alloc;
var param_7 : u32;
var raw1 : u32;
var param_8 : Alloc;
var param_9 : u32;
var raw2 : u32;
var param_10 : Alloc;
var param_11 : u32;
var s_17 : Path;
let x_355 : u32 = (*(ref_2)).offset;
ix_1 = (x_355 >> bitcast<u32>(2i));
let x_358 : u32 = ix_1;
let x_361 : Alloc = *(a_2);
param_6 = x_361;
param_7 = (x_358 + 0u);
let x_363 : u32 = read_mem_struct_Alloc_u11_u1_(&(param_6), &(param_7));
raw0_1 = x_363;
let x_365 : u32 = ix_1;
let x_368 : Alloc = *(a_2);
param_8 = x_368;
param_9 = (x_365 + 1u);
let x_370 : u32 = read_mem_struct_Alloc_u11_u1_(&(param_8), &(param_9));
raw1 = x_370;
let x_372 : u32 = ix_1;
let x_376 : Alloc = *(a_2);
param_10 = x_376;
param_11 = (x_372 + 2u);
let x_378 : u32 = read_mem_struct_Alloc_u11_u1_(&(param_10), &(param_11));
raw2 = x_378;
let x_381 : u32 = raw0_1;
let x_384 : u32 = raw0_1;
let x_387 : u32 = raw1;
let x_389 : u32 = raw1;
s_17.bbox = vec4<u32>((x_381 & 65535u), (x_384 >> bitcast<u32>(16i)), (x_387 & 65535u), (x_389 >> bitcast<u32>(16i)));
let x_395 : u32 = raw2;
s_17.tiles = TileRef(x_395);
let x_398 : Path = s_17;
return x_398;
}
fn write_tile_alloc_u1_struct_Alloc_u11_(el_ix : ptr<function, u32>, a_23 : ptr<function, Alloc>) {
return;
}
fn read_tile_alloc_u1_b1_(el_ix_1 : ptr<function, u32>, mem_ok_1 : ptr<function, bool>) -> Alloc {
var param_145 : u32;
var param_146 : u32;
var param_147 : bool;
param_145 = 0u;
let x_887 : u32 = x_883.conf.mem_size;
param_146 = x_887;
let x_889 : bool = *(mem_ok_1);
param_147 = x_889;
let x_890 : Alloc = new_alloc_u1_u1_b1_(&(param_145), &(param_146), &(param_147));
return x_890;
}
fn Tile_read_struct_Alloc_u11_struct_TileRef_u11_(a_3 : ptr<function, Alloc>, ref_3 : ptr<function, TileRef>) -> Tile {
var ix_2 : u32;
var raw0_2 : u32;
var param_12 : Alloc;
var param_13 : u32;
var raw1_1 : u32;
var param_14 : Alloc;
var param_15 : u32;
var s_18 : Tile;
let x_403 : u32 = (*(ref_3)).offset;
ix_2 = (x_403 >> bitcast<u32>(2i));
let x_406 : u32 = ix_2;
let x_409 : Alloc = *(a_3);
param_12 = x_409;
param_13 = (x_406 + 0u);
let x_411 : u32 = read_mem_struct_Alloc_u11_u1_(&(param_12), &(param_13));
raw0_2 = x_411;
let x_413 : u32 = ix_2;
let x_416 : Alloc = *(a_3);
param_14 = x_416;
param_15 = (x_413 + 1u);
let x_418 : u32 = read_mem_struct_Alloc_u11_u1_(&(param_14), &(param_15));
raw1_1 = x_418;
let x_420 : u32 = raw0_2;
s_18.tile = TileSegRef(x_420);
let x_424 : u32 = raw1_1;
s_18.backdrop = bitcast<i32>(x_424);
let x_428 : Tile = s_18;
return x_428;
}
fn malloc_stage_u1_u1_u1_(size_1 : ptr<function, u32>, mem_size : ptr<function, u32>, stage : ptr<function, u32>) -> u32 {
var offset_6 : u32;
let x_273 : u32 = *(size_1);
let x_276 : u32 = 0u;
offset_6 = x_276;
let x_277 : u32 = offset_6;
let x_278 : u32 = *(size_1);
let x_280 : u32 = *(mem_size);
if (((x_277 + x_278) > x_280)) {
offset_6 = 0u;
}
let x_284 : u32 = offset_6;
return x_284;
}
fn write_mem_struct_Alloc_u11_u1_u1_(alloc_1 : ptr<function, Alloc>, offset_3 : ptr<function, u32>, val : ptr<function, u32>) {
var param : Alloc;
var param_1 : u32;
let x_290 : Alloc = *(alloc_1);
param = x_290;
let x_292 : u32 = *(offset_3);
param_1 = x_292;
let x_293 : bool = touch_mem_struct_Alloc_u11_u1_(&(param), &(param_1));
if (!(x_293)) {
return;
}
let x_299 : u32 = *(offset_3);
let x_300 : u32 = *(val);
x_270.memory[x_299] = x_300;
return;
}
fn CmdJump_write_struct_Alloc_u11_struct_CmdJumpRef_u11_struct_CmdJump_u11_(a_11 : ptr<function, Alloc>, ref_11 : ptr<function, CmdJumpRef>, s_7 : ptr<function, CmdJump>) {
var ix_10 : u32;
var param_85 : Alloc;
var param_86 : u32;
var param_87 : u32;
let x_699 : u32 = (*(ref_11)).offset;
ix_10 = (x_699 >> bitcast<u32>(2i));
let x_701 : u32 = ix_10;
let x_704 : Alloc = *(a_11);
param_85 = x_704;
param_86 = (x_701 + 0u);
let x_708 : u32 = (*(s_7)).new_ref;
param_87 = x_708;
write_mem_struct_Alloc_u11_u1_u1_(&(param_85), &(param_86), &(param_87));
return;
}
fn Cmd_Jump_write_struct_Alloc_u11_struct_CmdRef_u11_struct_CmdJump_u11_(a_22 : ptr<function, Alloc>, ref_22 : ptr<function, CmdRef>, s_15 : ptr<function, CmdJump>) {
var param_139 : Alloc;
var param_140 : u32;
var param_141 : u32;
var param_142 : Alloc;
var param_143 : CmdJumpRef;
var param_144 : CmdJump;
let x_861 : u32 = (*(ref_22)).offset;
let x_865 : Alloc = *(a_22);
param_139 = x_865;
param_140 = (x_861 >> bitcast<u32>(2i));
param_141 = 11u;
write_mem_struct_Alloc_u11_u1_u1_(&(param_139), &(param_140), &(param_141));
let x_870 : u32 = (*(ref_22)).offset;
let x_874 : Alloc = *(a_22);
param_142 = x_874;
param_143 = CmdJumpRef((x_870 + 4u));
let x_877 : CmdJump = *(s_15);
param_144 = x_877;
CmdJump_write_struct_Alloc_u11_struct_CmdJumpRef_u11_struct_CmdJump_u11_(&(param_142), &(param_143), &(param_144));
return;
}
fn alloc_cmd_struct_Alloc_u11_struct_CmdRef_u11_u1_(cmd_alloc : ptr<function, Alloc>, cmd_ref : ptr<function, CmdRef>, cmd_limit : ptr<function, u32>) {
var new_cmd : u32;
var param_148 : u32;
var param_149 : u32;
var param_150 : u32;
var jump : CmdJump;
var param_151 : Alloc;
var param_152 : CmdRef;
var param_153 : CmdJump;
var param_154 : u32;
var param_155 : u32;
var param_156 : bool;
let x_894 : u32 = (*(cmd_ref)).offset;
let x_895 : u32 = *(cmd_limit);
if ((x_894 < x_895)) {
return;
}
param_148 = 1024u;
let x_905 : u32 = x_883.conf.mem_size;
param_149 = x_905;
param_150 = 8u;
let x_907 : u32 = malloc_stage_u1_u1_u1_(&(param_148), &(param_149), &(param_150));
new_cmd = x_907;
let x_908 : u32 = new_cmd;
if ((x_908 == 0u)) {
mem_ok_2 = false;
}
let x_913 : bool = mem_ok_2;
if (x_913) {
let x_917 : u32 = new_cmd;
jump = CmdJump(x_917);
let x_920 : Alloc = *(cmd_alloc);
param_151 = x_920;
let x_922 : CmdRef = *(cmd_ref);
param_152 = x_922;
let x_924 : CmdJump = jump;
param_153 = x_924;
Cmd_Jump_write_struct_Alloc_u11_struct_CmdRef_u11_struct_CmdJump_u11_(&(param_151), &(param_152), &(param_153));
}
let x_927 : u32 = new_cmd;
param_154 = x_927;
param_155 = 1024u;
param_156 = true;
let x_930 : Alloc = new_alloc_u1_u1_b1_(&(param_154), &(param_155), &(param_156));
*(cmd_alloc) = x_930;
let x_931 : u32 = new_cmd;
*(cmd_ref) = CmdRef(x_931);
let x_933 : u32 = new_cmd;
*(cmd_limit) = ((x_933 + 1024u) - 144u);
return;
}
fn CmdFill_write_struct_Alloc_u11_struct_CmdFillRef_u11_struct_CmdFill_u1_i11_(a_5 : ptr<function, Alloc>, ref_5 : ptr<function, CmdFillRef>, s_1 : ptr<function, CmdFill>) {
var ix_4 : u32;
var param_22 : Alloc;
var param_23 : u32;
var param_24 : u32;
var param_25 : Alloc;
var param_26 : u32;
var param_27 : u32;
let x_456 : u32 = (*(ref_5)).offset;
ix_4 = (x_456 >> bitcast<u32>(2i));
let x_458 : u32 = ix_4;
let x_461 : Alloc = *(a_5);
param_22 = x_461;
param_23 = (x_458 + 0u);
let x_465 : u32 = (*(s_1)).tile_ref;
param_24 = x_465;
write_mem_struct_Alloc_u11_u1_u1_(&(param_22), &(param_23), &(param_24));
let x_467 : u32 = ix_4;
let x_470 : i32 = (*(s_1)).backdrop;
let x_473 : Alloc = *(a_5);
param_25 = x_473;
param_26 = (x_467 + 1u);
param_27 = bitcast<u32>(x_470);
write_mem_struct_Alloc_u11_u1_u1_(&(param_25), &(param_26), &(param_27));
return;
}
fn Cmd_Fill_write_struct_Alloc_u11_struct_CmdRef_u11_struct_CmdFill_u1_i11_(a_13 : ptr<function, Alloc>, ref_13 : ptr<function, CmdRef>, s_8 : ptr<function, CmdFill>) {
var param_91 : Alloc;
var param_92 : u32;
var param_93 : u32;
var param_94 : Alloc;
var param_95 : CmdFillRef;
var param_96 : CmdFill;
let x_719 : u32 = (*(ref_13)).offset;
let x_722 : Alloc = *(a_13);
param_91 = x_722;
param_92 = (x_719 >> bitcast<u32>(2i));
param_93 = 1u;
write_mem_struct_Alloc_u11_u1_u1_(&(param_91), &(param_92), &(param_93));
let x_727 : u32 = (*(ref_13)).offset;
let x_731 : Alloc = *(a_13);
param_94 = x_731;
param_95 = CmdFillRef((x_727 + 4u));
let x_734 : CmdFill = *(s_8);
param_96 = x_734;
CmdFill_write_struct_Alloc_u11_struct_CmdFillRef_u11_struct_CmdFill_u1_i11_(&(param_94), &(param_95), &(param_96));
return;
}
fn Cmd_Solid_write_struct_Alloc_u11_struct_CmdRef_u11_(a_15 : ptr<function, Alloc>, ref_15 : ptr<function, CmdRef>) {
var param_103 : Alloc;
var param_104 : u32;
var param_105 : u32;
let x_755 : u32 = (*(ref_15)).offset;
let x_758 : Alloc = *(a_15);
param_103 = x_758;
param_104 = (x_755 >> bitcast<u32>(2i));
param_105 = 3u;
write_mem_struct_Alloc_u11_u1_u1_(&(param_103), &(param_104), &(param_105));
return;
}
fn CmdStroke_write_struct_Alloc_u11_struct_CmdStrokeRef_u11_struct_CmdStroke_u1_f11_(a_4 : ptr<function, Alloc>, ref_4 : ptr<function, CmdStrokeRef>, s : ptr<function, CmdStroke>) {
var ix_3 : u32;
var param_16 : Alloc;
var param_17 : u32;
var param_18 : u32;
var param_19 : Alloc;
var param_20 : u32;
var param_21 : u32;
let x_433 : u32 = (*(ref_4)).offset;
ix_3 = (x_433 >> bitcast<u32>(2i));
let x_435 : u32 = ix_3;
let x_438 : Alloc = *(a_4);
param_16 = x_438;
param_17 = (x_435 + 0u);
let x_442 : u32 = (*(s)).tile_ref;
param_18 = x_442;
write_mem_struct_Alloc_u11_u1_u1_(&(param_16), &(param_17), &(param_18));
let x_444 : u32 = ix_3;
let x_447 : f32 = (*(s)).half_width;
let x_450 : Alloc = *(a_4);
param_19 = x_450;
param_20 = (x_444 + 1u);
param_21 = bitcast<u32>(x_447);
write_mem_struct_Alloc_u11_u1_u1_(&(param_19), &(param_20), &(param_21));
return;
}
fn Cmd_Stroke_write_struct_Alloc_u11_struct_CmdRef_u11_struct_CmdStroke_u1_f11_(a_14 : ptr<function, Alloc>, ref_14 : ptr<function, CmdRef>, s_9 : ptr<function, CmdStroke>) {
var param_97 : Alloc;
var param_98 : u32;
var param_99 : u32;
var param_100 : Alloc;
var param_101 : CmdStrokeRef;
var param_102 : CmdStroke;
let x_737 : u32 = (*(ref_14)).offset;
let x_740 : Alloc = *(a_14);
param_97 = x_740;
param_98 = (x_737 >> bitcast<u32>(2i));
param_99 = 2u;
write_mem_struct_Alloc_u11_u1_u1_(&(param_97), &(param_98), &(param_99));
let x_745 : u32 = (*(ref_14)).offset;
let x_749 : Alloc = *(a_14);
param_100 = x_749;
param_101 = CmdStrokeRef((x_745 + 4u));
let x_752 : CmdStroke = *(s_9);
param_102 = x_752;
CmdStroke_write_struct_Alloc_u11_struct_CmdStrokeRef_u11_struct_CmdStroke_u1_f11_(&(param_100), &(param_101), &(param_102));
return;
}
fn write_fill_struct_Alloc_u11_struct_CmdRef_u11_struct_Tile_struct_TileSegRef_u11_i11_f1_(alloc_3 : ptr<function, Alloc>, cmd_ref_1 : ptr<function, CmdRef>, tile : ptr<function, Tile>, linewidth : ptr<function, f32>) {
var cmd_fill : CmdFill;
var param_157 : Alloc;
var param_158 : CmdRef;
var param_159 : CmdFill;
var param_160 : Alloc;
var param_161 : CmdRef;
var cmd_stroke : CmdStroke;
var param_162 : Alloc;
var param_163 : CmdRef;
var param_164 : CmdStroke;
let x_937 : f32 = *(linewidth);
if ((x_937 < 0.0f)) {
let x_943 : u32 = (*(tile)).tile.offset;
if ((x_943 != 0u)) {
let x_949 : u32 = (*(tile)).tile.offset;
let x_951 : i32 = (*(tile)).backdrop;
cmd_fill = CmdFill(x_949, x_951);
let x_953 : bool = mem_ok_2;
if (x_953) {
let x_957 : Alloc = *(alloc_3);
param_157 = x_957;
let x_959 : CmdRef = *(cmd_ref_1);
param_158 = x_959;
let x_961 : CmdFill = cmd_fill;
param_159 = x_961;
Cmd_Fill_write_struct_Alloc_u11_struct_CmdRef_u11_struct_CmdFill_u1_i11_(&(param_157), &(param_158), &(param_159));
}
let x_965 : u32 = (*(cmd_ref_1)).offset;
(*(cmd_ref_1)).offset = (x_965 + 12u);
} else {
let x_969 : bool = mem_ok_2;
if (x_969) {
let x_973 : Alloc = *(alloc_3);
param_160 = x_973;
let x_975 : CmdRef = *(cmd_ref_1);
param_161 = x_975;
Cmd_Solid_write_struct_Alloc_u11_struct_CmdRef_u11_(&(param_160), &(param_161));
}
let x_978 : u32 = (*(cmd_ref_1)).offset;
(*(cmd_ref_1)).offset = (x_978 + 4u);
}
} else {
let x_984 : u32 = (*(tile)).tile.offset;
let x_986 : f32 = *(linewidth);
cmd_stroke = CmdStroke(x_984, (0.5f * x_986));
let x_989 : bool = mem_ok_2;
if (x_989) {
let x_993 : Alloc = *(alloc_3);
param_162 = x_993;
let x_995 : CmdRef = *(cmd_ref_1);
param_163 = x_995;
let x_997 : CmdStroke = cmd_stroke;
param_164 = x_997;
Cmd_Stroke_write_struct_Alloc_u11_struct_CmdRef_u11_struct_CmdStroke_u1_f11_(&(param_162), &(param_163), &(param_164));
}
let x_1000 : u32 = (*(cmd_ref_1)).offset;
(*(cmd_ref_1)).offset = (x_1000 + 12u);
}
return;
}
fn CmdColor_write_struct_Alloc_u11_struct_CmdColorRef_u11_struct_CmdColor_u11_(a_6 : ptr<function, Alloc>, ref_6 : ptr<function, CmdColorRef>, s_2 : ptr<function, CmdColor>) {
var ix_5 : u32;
var param_28 : Alloc;
var param_29 : u32;
var param_30 : u32;
let x_479 : u32 = (*(ref_6)).offset;
ix_5 = (x_479 >> bitcast<u32>(2i));
let x_481 : u32 = ix_5;
let x_484 : Alloc = *(a_6);
param_28 = x_484;
param_29 = (x_481 + 0u);
let x_488 : u32 = (*(s_2)).rgba_color;
param_30 = x_488;
write_mem_struct_Alloc_u11_u1_u1_(&(param_28), &(param_29), &(param_30));
return;
}
fn Cmd_Color_write_struct_Alloc_u11_struct_CmdRef_u11_struct_CmdColor_u11_(a_16 : ptr<function, Alloc>, ref_16 : ptr<function, CmdRef>, s_10 : ptr<function, CmdColor>) {
var param_106 : Alloc;
var param_107 : u32;
var param_108 : u32;
var param_109 : Alloc;
var param_110 : CmdColorRef;
var param_111 : CmdColor;
let x_763 : u32 = (*(ref_16)).offset;
let x_766 : Alloc = *(a_16);
param_106 = x_766;
param_107 = (x_763 >> bitcast<u32>(2i));
param_108 = 5u;
write_mem_struct_Alloc_u11_u1_u1_(&(param_106), &(param_107), &(param_108));
let x_771 : u32 = (*(ref_16)).offset;
let x_775 : Alloc = *(a_16);
param_109 = x_775;
param_110 = CmdColorRef((x_771 + 4u));
let x_778 : CmdColor = *(s_10);
param_111 = x_778;
CmdColor_write_struct_Alloc_u11_struct_CmdColorRef_u11_struct_CmdColor_u11_(&(param_109), &(param_110), &(param_111));
return;
}
fn CmdLinGrad_write_struct_Alloc_u11_struct_CmdLinGradRef_u11_struct_CmdLinGrad_u1_f1_f1_f11_(a_7 : ptr<function, Alloc>, ref_7 : ptr<function, CmdLinGradRef>, s_3 : ptr<function, CmdLinGrad>) {
var ix_6 : u32;
var param_31 : Alloc;
var param_32 : u32;
var param_33 : u32;
var param_34 : Alloc;
var param_35 : u32;
var param_36 : u32;
var param_37 : Alloc;
var param_38 : u32;
var param_39 : u32;
var param_40 : Alloc;
var param_41 : u32;
var param_42 : u32;
let x_492 : u32 = (*(ref_7)).offset;
ix_6 = (x_492 >> bitcast<u32>(2i));
let x_494 : u32 = ix_6;
let x_497 : Alloc = *(a_7);
param_31 = x_497;
param_32 = (x_494 + 0u);
let x_501 : u32 = (*(s_3)).index;
param_33 = x_501;
write_mem_struct_Alloc_u11_u1_u1_(&(param_31), &(param_32), &(param_33));
let x_503 : u32 = ix_6;
let x_506 : f32 = (*(s_3)).line_x;
let x_509 : Alloc = *(a_7);
param_34 = x_509;
param_35 = (x_503 + 1u);
param_36 = bitcast<u32>(x_506);
write_mem_struct_Alloc_u11_u1_u1_(&(param_34), &(param_35), &(param_36));
let x_513 : u32 = ix_6;
let x_516 : f32 = (*(s_3)).line_y;
let x_519 : Alloc = *(a_7);
param_37 = x_519;
param_38 = (x_513 + 2u);
param_39 = bitcast<u32>(x_516);
write_mem_struct_Alloc_u11_u1_u1_(&(param_37), &(param_38), &(param_39));
let x_523 : u32 = ix_6;
let x_527 : f32 = (*(s_3)).line_c;
let x_530 : Alloc = *(a_7);
param_40 = x_530;
param_41 = (x_523 + 3u);
param_42 = bitcast<u32>(x_527);
write_mem_struct_Alloc_u11_u1_u1_(&(param_40), &(param_41), &(param_42));
return;
}
fn Cmd_LinGrad_write_struct_Alloc_u11_struct_CmdRef_u11_struct_CmdLinGrad_u1_f1_f1_f11_(a_17 : ptr<function, Alloc>, ref_17 : ptr<function, CmdRef>, s_11 : ptr<function, CmdLinGrad>) {
var param_112 : Alloc;
var param_113 : u32;
var param_114 : u32;
var param_115 : Alloc;
var param_116 : CmdLinGradRef;
var param_117 : CmdLinGrad;
let x_781 : u32 = (*(ref_17)).offset;
let x_784 : Alloc = *(a_17);
param_112 = x_784;
param_113 = (x_781 >> bitcast<u32>(2i));
param_114 = 6u;
write_mem_struct_Alloc_u11_u1_u1_(&(param_112), &(param_113), &(param_114));
let x_789 : u32 = (*(ref_17)).offset;
let x_793 : Alloc = *(a_17);
param_115 = x_793;
param_116 = CmdLinGradRef((x_789 + 4u));
let x_796 : CmdLinGrad = *(s_11);
param_117 = x_796;
CmdLinGrad_write_struct_Alloc_u11_struct_CmdLinGradRef_u11_struct_CmdLinGrad_u1_f1_f1_f11_(&(param_115), &(param_116), &(param_117));
return;
}
fn CmdRadGrad_write_struct_Alloc_u11_struct_CmdRadGradRef_u11_struct_CmdRadGrad_u1_vf4_vf2_vf2_f1_f11_(a_8 : ptr<function, Alloc>, ref_8 : ptr<function, CmdRadGradRef>, s_4 : ptr<function, CmdRadGrad>) {
var ix_7 : u32;
var param_43 : Alloc;
var param_44 : u32;
var param_45 : u32;
var param_46 : Alloc;
var param_47 : u32;
var param_48 : u32;
var param_49 : Alloc;
var param_50 : u32;
var param_51 : u32;
var param_52 : Alloc;
var param_53 : u32;
var param_54 : u32;
var param_55 : Alloc;
var param_56 : u32;
var param_57 : u32;
var param_58 : Alloc;
var param_59 : u32;
var param_60 : u32;
var param_61 : Alloc;
var param_62 : u32;
var param_63 : u32;
var param_64 : Alloc;
var param_65 : u32;
var param_66 : u32;
var param_67 : Alloc;
var param_68 : u32;
var param_69 : u32;
var param_70 : Alloc;
var param_71 : u32;
var param_72 : u32;
var param_73 : Alloc;
var param_74 : u32;
var param_75 : u32;
let x_536 : u32 = (*(ref_8)).offset;
ix_7 = (x_536 >> bitcast<u32>(2i));
let x_538 : u32 = ix_7;
let x_541 : Alloc = *(a_8);
param_43 = x_541;
param_44 = (x_538 + 0u);
let x_545 : u32 = (*(s_4)).index;
param_45 = x_545;
write_mem_struct_Alloc_u11_u1_u1_(&(param_43), &(param_44), &(param_45));
let x_547 : u32 = ix_7;
let x_550 : f32 = (*(s_4)).mat_.x;
let x_553 : Alloc = *(a_8);
param_46 = x_553;
param_47 = (x_547 + 1u);
param_48 = bitcast<u32>(x_550);
write_mem_struct_Alloc_u11_u1_u1_(&(param_46), &(param_47), &(param_48));
let x_557 : u32 = ix_7;
let x_560 : f32 = (*(s_4)).mat_.y;
let x_563 : Alloc = *(a_8);
param_49 = x_563;
param_50 = (x_557 + 2u);
param_51 = bitcast<u32>(x_560);
write_mem_struct_Alloc_u11_u1_u1_(&(param_49), &(param_50), &(param_51));
let x_567 : u32 = ix_7;
let x_570 : f32 = (*(s_4)).mat_.z;
let x_573 : Alloc = *(a_8);
param_52 = x_573;
param_53 = (x_567 + 3u);
param_54 = bitcast<u32>(x_570);
write_mem_struct_Alloc_u11_u1_u1_(&(param_52), &(param_53), &(param_54));
let x_577 : u32 = ix_7;
let x_580 : f32 = (*(s_4)).mat_.w;
let x_583 : Alloc = *(a_8);
param_55 = x_583;
param_56 = (x_577 + 4u);
param_57 = bitcast<u32>(x_580);
write_mem_struct_Alloc_u11_u1_u1_(&(param_55), &(param_56), &(param_57));
let x_587 : u32 = ix_7;
let x_591 : f32 = (*(s_4)).xlat.x;
let x_594 : Alloc = *(a_8);
param_58 = x_594;
param_59 = (x_587 + 5u);
param_60 = bitcast<u32>(x_591);
write_mem_struct_Alloc_u11_u1_u1_(&(param_58), &(param_59), &(param_60));
let x_598 : u32 = ix_7;
let x_602 : f32 = (*(s_4)).xlat.y;
let x_605 : Alloc = *(a_8);
param_61 = x_605;
param_62 = (x_598 + 6u);
param_63 = bitcast<u32>(x_602);
write_mem_struct_Alloc_u11_u1_u1_(&(param_61), &(param_62), &(param_63));
let x_609 : u32 = ix_7;
let x_613 : f32 = (*(s_4)).c1.x;
let x_616 : Alloc = *(a_8);
param_64 = x_616;
param_65 = (x_609 + 7u);
param_66 = bitcast<u32>(x_613);
write_mem_struct_Alloc_u11_u1_u1_(&(param_64), &(param_65), &(param_66));
let x_620 : u32 = ix_7;
let x_624 : f32 = (*(s_4)).c1.y;
let x_627 : Alloc = *(a_8);
param_67 = x_627;
param_68 = (x_620 + 8u);
param_69 = bitcast<u32>(x_624);
write_mem_struct_Alloc_u11_u1_u1_(&(param_67), &(param_68), &(param_69));
let x_631 : u32 = ix_7;
let x_636 : f32 = (*(s_4)).ra;
let x_639 : Alloc = *(a_8);
param_70 = x_639;
param_71 = (x_631 + 9u);
param_72 = bitcast<u32>(x_636);
write_mem_struct_Alloc_u11_u1_u1_(&(param_70), &(param_71), &(param_72));
let x_643 : u32 = ix_7;
let x_648 : f32 = (*(s_4)).roff;
let x_651 : Alloc = *(a_8);
param_73 = x_651;
param_74 = (x_643 + 10u);
param_75 = bitcast<u32>(x_648);
write_mem_struct_Alloc_u11_u1_u1_(&(param_73), &(param_74), &(param_75));
return;
}
fn Cmd_RadGrad_write_struct_Alloc_u11_struct_CmdRef_u11_struct_CmdRadGrad_u1_vf4_vf2_vf2_f1_f11_(a_18 : ptr<function, Alloc>, ref_18 : ptr<function, CmdRef>, s_12 : ptr<function, CmdRadGrad>) {
var param_118 : Alloc;
var param_119 : u32;
var param_120 : u32;
var param_121 : Alloc;
var param_122 : CmdRadGradRef;
var param_123 : CmdRadGrad;
let x_799 : u32 = (*(ref_18)).offset;
let x_802 : Alloc = *(a_18);
param_118 = x_802;
param_119 = (x_799 >> bitcast<u32>(2i));
param_120 = 7u;
write_mem_struct_Alloc_u11_u1_u1_(&(param_118), &(param_119), &(param_120));
let x_807 : u32 = (*(ref_18)).offset;
let x_811 : Alloc = *(a_18);
param_121 = x_811;
param_122 = CmdRadGradRef((x_807 + 4u));
let x_814 : CmdRadGrad = *(s_12);
param_123 = x_814;
CmdRadGrad_write_struct_Alloc_u11_struct_CmdRadGradRef_u11_struct_CmdRadGrad_u1_vf4_vf2_vf2_f1_f11_(&(param_121), &(param_122), &(param_123));
return;
}
fn CmdImage_write_struct_Alloc_u11_struct_CmdImageRef_u11_struct_CmdImage_u1_vi21_(a_9 : ptr<function, Alloc>, ref_9 : ptr<function, CmdImageRef>, s_5 : ptr<function, CmdImage>) {
var ix_8 : u32;
var param_76 : Alloc;
var param_77 : u32;
var param_78 : u32;
var param_79 : Alloc;
var param_80 : u32;
var param_81 : u32;
let x_657 : u32 = (*(ref_9)).offset;
ix_8 = (x_657 >> bitcast<u32>(2i));
let x_659 : u32 = ix_8;
let x_662 : Alloc = *(a_9);
param_76 = x_662;
param_77 = (x_659 + 0u);
let x_666 : u32 = (*(s_5)).index;
param_78 = x_666;
write_mem_struct_Alloc_u11_u1_u1_(&(param_76), &(param_77), &(param_78));
let x_668 : u32 = ix_8;
let x_671 : i32 = (*(s_5)).offset.x;
let x_675 : i32 = (*(s_5)).offset.y;
let x_680 : Alloc = *(a_9);
param_79 = x_680;
param_80 = (x_668 + 1u);
param_81 = ((bitcast<u32>(x_671) & 65535u) | (bitcast<u32>(x_675) << bitcast<u32>(16i)));
write_mem_struct_Alloc_u11_u1_u1_(&(param_79), &(param_80), &(param_81));
return;
}
fn Cmd_Image_write_struct_Alloc_u11_struct_CmdRef_u11_struct_CmdImage_u1_vi21_(a_19 : ptr<function, Alloc>, ref_19 : ptr<function, CmdRef>, s_13 : ptr<function, CmdImage>) {
var param_124 : Alloc;
var param_125 : u32;
var param_126 : u32;
var param_127 : Alloc;
var param_128 : CmdImageRef;
var param_129 : CmdImage;
let x_817 : u32 = (*(ref_19)).offset;
let x_820 : Alloc = *(a_19);
param_124 = x_820;
param_125 = (x_817 >> bitcast<u32>(2i));
param_126 = 8u;
write_mem_struct_Alloc_u11_u1_u1_(&(param_124), &(param_125), &(param_126));
let x_825 : u32 = (*(ref_19)).offset;
let x_829 : Alloc = *(a_19);
param_127 = x_829;
param_128 = CmdImageRef((x_825 + 4u));
let x_832 : CmdImage = *(s_13);
param_129 = x_832;
CmdImage_write_struct_Alloc_u11_struct_CmdImageRef_u11_struct_CmdImage_u1_vi21_(&(param_127), &(param_128), &(param_129));
return;
}
fn Cmd_BeginClip_write_struct_Alloc_u11_struct_CmdRef_u11_(a_20 : ptr<function, Alloc>, ref_20 : ptr<function, CmdRef>) {
var param_130 : Alloc;
var param_131 : u32;
var param_132 : u32;
let x_835 : u32 = (*(ref_20)).offset;
let x_838 : Alloc = *(a_20);
param_130 = x_838;
param_131 = (x_835 >> bitcast<u32>(2i));
param_132 = 9u;
write_mem_struct_Alloc_u11_u1_u1_(&(param_130), &(param_131), &(param_132));
return;
}
fn CmdEndClip_write_struct_Alloc_u11_struct_CmdEndClipRef_u11_struct_CmdEndClip_u11_(a_10 : ptr<function, Alloc>, ref_10 : ptr<function, CmdEndClipRef>, s_6 : ptr<function, CmdEndClip>) {
var ix_9 : u32;
var param_82 : Alloc;
var param_83 : u32;
var param_84 : u32;
let x_686 : u32 = (*(ref_10)).offset;
ix_9 = (x_686 >> bitcast<u32>(2i));
let x_688 : u32 = ix_9;
let x_691 : Alloc = *(a_10);
param_82 = x_691;
param_83 = (x_688 + 0u);
let x_695 : u32 = (*(s_6)).blend;
param_84 = x_695;
write_mem_struct_Alloc_u11_u1_u1_(&(param_82), &(param_83), &(param_84));
return;
}
fn Cmd_EndClip_write_struct_Alloc_u11_struct_CmdRef_u11_struct_CmdEndClip_u11_(a_21 : ptr<function, Alloc>, ref_21 : ptr<function, CmdRef>, s_14 : ptr<function, CmdEndClip>) {
var param_133 : Alloc;
var param_134 : u32;
var param_135 : u32;
var param_136 : Alloc;
var param_137 : CmdEndClipRef;
var param_138 : CmdEndClip;
let x_843 : u32 = (*(ref_21)).offset;
let x_846 : Alloc = *(a_21);
param_133 = x_846;
param_134 = (x_843 >> bitcast<u32>(2i));
param_135 = 10u;
write_mem_struct_Alloc_u11_u1_u1_(&(param_133), &(param_134), &(param_135));
let x_851 : u32 = (*(ref_21)).offset;
let x_855 : Alloc = *(a_21);
param_136 = x_855;
param_137 = CmdEndClipRef((x_851 + 4u));
let x_858 : CmdEndClip = *(s_14);
param_138 = x_858;
CmdEndClip_write_struct_Alloc_u11_struct_CmdEndClipRef_u11_struct_CmdEndClip_u11_(&(param_136), &(param_137), &(param_138));
return;
}
fn Cmd_End_write_struct_Alloc_u11_struct_CmdRef_u11_(a_12 : ptr<function, Alloc>, ref_12 : ptr<function, CmdRef>) {
var param_88 : Alloc;
var param_89 : u32;
var param_90 : u32;
let x_711 : u32 = (*(ref_12)).offset;
let x_714 : Alloc = *(a_12);
param_88 = x_714;
param_89 = (x_711 >> bitcast<u32>(2i));
param_90 = 0u;
write_mem_struct_Alloc_u11_u1_u1_(&(param_88), &(param_89), &(param_90));
return;
}
fn main_1() {
var param_165 : u32;
var width_in_bins : u32;
var bin_ix : u32;
var partition_ix : u32;
var n_partitions : u32;
var th_ix : u32;
var bin_tile_x : u32;
var bin_tile_y : u32;
var tile_x : u32;
var tile_y : u32;
var this_tile_ix : u32;
var cmd_alloc_1 : Alloc;
var param_166 : Alloc;
var param_167 : u32;
var param_168 : u32;
var cmd_ref_2 : CmdRef;
var cmd_limit_1 : u32;
var clip_depth : u32;
var clip_zero_depth : u32;
var rd_ix : u32;
var wr_ix : u32;
var part_start_ix : u32;
var ready_ix : u32;
var scratch_alloc : Alloc;
var param_169 : Alloc;
var param_170 : u32;
var param_171 : u32;
var render_blend_depth : u32;
var max_blend_depth : u32;
var drawmonoid_start : u32;
var drawtag_start : u32;
var drawdata_start : u32;
var drawinfo_start : u32;
var i : u32;
var count : u32;
var in_ix : u32;
var param_172 : Alloc;
var param_173 : u32;
var offset_7 : u32;
var param_174 : Alloc;
var param_175 : u32;
var param_176 : u32;
var param_177 : u32;
var param_178 : bool;
var i_1 : u32;
var ix_11 : u32;
var part_ix : u32;
var i_2 : u32;
var probe : u32;
var x_1314 : u32;
var bin_alloc : Alloc;
var inst_ref : BinInstanceRef;
var inst : BinInstance;
var param_179 : BinInstanceRef;
var param_180 : u32;
var param_181 : Alloc;
var param_182 : BinInstanceRef;
var tag : u32;
var element_ix : u32;
var drawmonoid_base : u32;
var path_ix : u32;
var path : Path;
var param_183 : Alloc;
var param_184 : PathRef;
var stride : u32;
var dx : i32;
var dy : i32;
var x0 : i32;
var y0 : i32;
var x1 : i32;
var y1 : i32;
var tile_count : u32;
var base : u32;
var path_alloc : Alloc;
var param_185 : u32;
var param_186 : u32;
var param_187 : bool;
var param_188 : u32;
var param_189 : Alloc;
var i_3 : u32;
var total_tile_count : u32;
var ix_12 : u32;
var el_ix_2 : u32;
var i_4 : u32;
var probe_1 : u32;
var element_ix_1 : u32;
var tag_1 : u32;
var seq_ix : u32;
var x_1614 : u32;
var width : u32;
var x : u32;
var y : u32;
var include_tile : bool;
var tile_1 : Tile;
var param_190 : u32;
var param_191 : bool;
var param_192 : Alloc;
var param_193 : TileRef;
var is_clip : bool;
var is_blend : bool;
var drawmonoid_base_1 : u32;
var scene_offset : u32;
var dd : u32;
var blend : u32;
var el_slice : u32;
var el_mask : u32;
var slice_ix : u32;
var bitmap : u32;
var element_ref_ix : u32;
var element_ix_2 : u32;
var drawtag : u32;
var tile_2 : Tile;
var param_194 : u32;
var param_195 : bool;
var param_196 : Alloc;
var param_197 : TileRef;
var drawmonoid_base_2 : u32;
var scene_offset_1 : u32;
var info_offset : u32;
var dd_1 : u32;
var di : u32;
var linewidth_1 : f32;
var param_198 : Alloc;
var param_199 : CmdRef;
var param_200 : u32;
var param_201 : Alloc;
var param_202 : CmdRef;
var param_203 : Tile;
var param_204 : f32;
var rgba : u32;
var param_205 : Alloc;
var param_206 : CmdRef;
var param_207 : CmdColor;
var param_208 : Alloc;
var param_209 : CmdRef;
var param_210 : u32;
var param_211 : Alloc;
var param_212 : CmdRef;
var param_213 : Tile;
var param_214 : f32;
var cmd_lin : CmdLinGrad;
var param_215 : Alloc;
var param_216 : CmdRef;
var param_217 : CmdLinGrad;
var param_218 : Alloc;
var param_219 : CmdRef;
var param_220 : u32;
var param_221 : Alloc;
var param_222 : CmdRef;
var param_223 : Tile;
var param_224 : f32;
var cmd_rad : CmdRadGrad;
var param_225 : Alloc;
var param_226 : CmdRef;
var param_227 : CmdRadGrad;
var param_228 : Alloc;
var param_229 : CmdRef;
var param_230 : u32;
var param_231 : Alloc;
var param_232 : CmdRef;
var param_233 : Tile;
var param_234 : f32;
var index_1 : u32;
var raw1_2 : u32;
var offset_8 : vec2<i32>;
var param_235 : Alloc;
var param_236 : CmdRef;
var param_237 : CmdImage;
var param_238 : Alloc;
var param_239 : CmdRef;
var param_240 : u32;
var param_241 : Alloc;
var param_242 : CmdRef;
var param_243 : Alloc;
var param_244 : CmdRef;
var param_245 : Tile;
var param_246 : f32;
var blend_1 : u32;
var param_247 : Alloc;
var param_248 : CmdRef;
var param_249 : CmdEndClip;
var param_250 : Alloc;
var param_251 : CmdRef;
var scratch_size : u32;
var scratch : u32;
var param_252 : Alloc;
var param_253 : u32;
var param_254 : u32;
var x_2219 : bool;
var x_2220 : bool;
mem_ok_2 = true;
param_165 = 7u;
let x_1004 : bool = check_deps_u1_(&(param_165));
if (!(x_1004)) {
return;
}
let x_1011 : u32 = x_883.conf.width_in_tiles;
width_in_bins = (((x_1011 + 16u) - 1u) / 16u);
let x_1017 : u32 = width_in_bins;
let x_1023 : u32 = gl_WorkGroupID.y;
let x_1026 : u32 = gl_WorkGroupID.x;
bin_ix = ((x_1017 * x_1023) + x_1026);
partition_ix = 0u;
let x_1031 : u32 = x_883.conf.n_elements;
n_partitions = (((x_1031 + 256u) - 1u) / 256u);
let x_1039 : u32 = gl_LocalInvocationID.x;
th_ix = x_1039;
let x_1042 : u32 = gl_WorkGroupID.x;
bin_tile_x = (16u * x_1042);
let x_1046 : u32 = gl_WorkGroupID.y;
bin_tile_y = (16u * x_1046);
let x_1050 : u32 = gl_LocalInvocationID.x;
tile_x = (x_1050 % 16u);
let x_1054 : u32 = gl_LocalInvocationID.x;
tile_y = (x_1054 / 16u);
let x_1057 : u32 = bin_tile_y;
let x_1058 : u32 = tile_y;
let x_1061 : u32 = x_883.conf.width_in_tiles;
let x_1063 : u32 = bin_tile_x;
let x_1065 : u32 = tile_x;
this_tile_ix = ((((x_1057 + x_1058) * x_1061) + x_1063) + x_1065);
let x_1069 : u32 = this_tile_ix;
let x_1074 : Alloc_1 = x_883.conf.ptcl_alloc;
param_166.offset = x_1074.offset;
param_167 = (x_1069 * 1024u);
param_168 = 1024u;
let x_1079 : Alloc = slice_mem_struct_Alloc_u11_u1_u1_(&(param_166), &(param_167), &(param_168));
cmd_alloc_1 = x_1079;
let x_1082 : u32 = cmd_alloc_1.offset;
cmd_ref_2 = CmdRef(x_1082);
let x_1086 : u32 = cmd_ref_2.offset;
cmd_limit_1 = ((x_1086 + 1024u) - 144u);
clip_depth = 0u;
clip_zero_depth = 0u;
rd_ix = 0u;
wr_ix = 0u;
part_start_ix = 0u;
ready_ix = 0u;
let x_1097 : Alloc = cmd_alloc_1;
param_169 = x_1097;
param_170 = 0u;
param_171 = 8u;
let x_1100 : Alloc = slice_mem_struct_Alloc_u11_u1_u1_(&(param_169), &(param_170), &(param_171));
scratch_alloc = x_1100;
let x_1102 : u32 = cmd_ref_2.offset;
cmd_ref_2.offset = (x_1102 + 4u);
render_blend_depth = 0u;
max_blend_depth = 0u;
let x_1110 : u32 = x_883.conf.drawmonoid_alloc.offset;
drawmonoid_start = (x_1110 >> bitcast<u32>(2i));
let x_1115 : u32 = x_883.conf.drawtag_offset;
drawtag_start = (x_1115 >> bitcast<u32>(2i));
let x_1120 : u32 = x_883.conf.drawdata_offset;
drawdata_start = (x_1120 >> bitcast<u32>(2i));
let x_1125 : u32 = x_883.conf.drawinfo_alloc.offset;
drawinfo_start = (x_1125 >> bitcast<u32>(2i));
loop {
if (true) {
} else {
break;
}
i = 0u;
loop {
let x_1138 : u32 = i;
if ((x_1138 < 8u)) {
} else {
break;
}
let x_1144 : u32 = i;
let x_1145 : u32 = th_ix;
sh_bitmaps[x_1144][x_1145] = 0u;
continuing {
let x_1148 : u32 = i;
i = (x_1148 + bitcast<u32>(1i));
}
}
loop {
var x_1173 : bool;
var x_1174 : bool;
let x_1154 : u32 = ready_ix;
let x_1155 : u32 = wr_ix;
let x_1157 : u32 = partition_ix;
let x_1158 : u32 = n_partitions;
if (((x_1154 == x_1155) & (x_1157 < x_1158))) {
let x_1163 : u32 = ready_ix;
part_start_ix = x_1163;
count = 0u;
let x_1165 : u32 = th_ix;
let x_1166 : bool = (x_1165 < 256u);
x_1174 = x_1166;
if (x_1166) {
let x_1169 : u32 = partition_ix;
let x_1170 : u32 = th_ix;
let x_1172 : u32 = n_partitions;
x_1173 = ((x_1169 + x_1170) < x_1172);
x_1174 = x_1173;
}
if (x_1174) {
let x_1180 : u32 = x_883.conf.bin_alloc.offset;
let x_1182 : u32 = partition_ix;
let x_1183 : u32 = th_ix;
let x_1186 : u32 = bin_ix;
in_ix = ((x_1180 >> bitcast<u32>(2i)) + ((((x_1182 + x_1183) * 256u) + x_1186) * 2u));
let x_1192 : Alloc_1 = x_883.conf.bin_alloc;
param_172.offset = x_1192.offset;
let x_1196 : u32 = in_ix;
param_173 = x_1196;
let x_1197 : u32 = read_mem_struct_Alloc_u11_u1_(&(param_172), &(param_173));
count = x_1197;
let x_1199 : u32 = in_ix;
let x_1203 : Alloc_1 = x_883.conf.bin_alloc;
param_174.offset = x_1203.offset;
param_175 = (x_1199 + 1u);
let x_1207 : u32 = read_mem_struct_Alloc_u11_u1_(&(param_174), &(param_175));
offset_7 = x_1207;
let x_1211 : u32 = th_ix;
let x_1212 : u32 = count;
let x_1215 : u32 = offset_7;
param_176 = x_1215;
param_177 = (x_1212 * 4u);
param_178 = true;
let x_1218 : Alloc = new_alloc_u1_u1_b1_(&(param_176), &(param_177), &(param_178));
sh_part_elements[x_1211] = x_1218;
}
i_1 = 0u;
loop {
let x_1227 : u32 = i_1;
if ((x_1227 < 8u)) {
} else {
break;
}
let x_1229 : u32 = th_ix;
if ((x_1229 < 256u)) {
let x_1235 : u32 = th_ix;
let x_1236 : u32 = count;
sh_part_count[x_1235] = x_1236;
}
workgroupBarrier();
let x_1239 : u32 = th_ix;
if ((x_1239 < 256u)) {
let x_1243 : u32 = th_ix;
let x_1244 : u32 = i_1;
if ((x_1243 >= (1u << x_1244))) {
let x_1249 : u32 = th_ix;
let x_1250 : u32 = i_1;
let x_1254 : u32 = sh_part_count[(x_1249 - (1u << x_1250))];
let x_1255 : u32 = count;
count = (x_1255 + x_1254);
}
}
workgroupBarrier();
continuing {
let x_1257 : u32 = i_1;
i_1 = (x_1257 + bitcast<u32>(1i));
}
}
let x_1259 : u32 = th_ix;
if ((x_1259 < 256u)) {
let x_1263 : u32 = th_ix;
let x_1264 : u32 = part_start_ix;
let x_1265 : u32 = count;
sh_part_count[x_1263] = (x_1264 + x_1265);
}
workgroupBarrier();
let x_1270 : u32 = sh_part_count[255i];
ready_ix = x_1270;
let x_1271 : u32 = partition_ix;
partition_ix = (x_1271 + 256u);
}
let x_1274 : u32 = rd_ix;
let x_1275 : u32 = th_ix;
ix_11 = (x_1274 + x_1275);
let x_1277 : u32 = ix_11;
let x_1278 : u32 = wr_ix;
let x_1280 : u32 = ix_11;
let x_1281 : u32 = ready_ix;
if (((x_1277 >= x_1278) & (x_1280 < x_1281))) {
part_ix = 0u;
i_2 = 0u;
loop {
let x_1293 : u32 = i_2;
if ((x_1293 < 8u)) {
} else {
break;
}
let x_1296 : u32 = part_ix;
let x_1298 : u32 = i_2;
probe = (x_1296 + (128u >> x_1298));
let x_1301 : u32 = ix_11;
let x_1302 : u32 = probe;
let x_1305 : u32 = sh_part_count[(x_1302 - 1u)];
if ((x_1301 >= x_1305)) {
let x_1309 : u32 = probe;
part_ix = x_1309;
}
continuing {
let x_1310 : u32 = i_2;
i_2 = (x_1310 + bitcast<u32>(1i));
}
}
let x_1312 : u32 = part_ix;
if ((x_1312 > 0u)) {
let x_1317 : u32 = part_ix;
let x_1320 : u32 = sh_part_count[(x_1317 - 1u)];
x_1314 = x_1320;
} else {
let x_1322 : u32 = part_start_ix;
x_1314 = x_1322;
}
let x_1323 : u32 = x_1314;
let x_1324 : u32 = ix_11;
ix_11 = (x_1324 - x_1323);
let x_1327 : u32 = part_ix;
let x_1329 : Alloc = sh_part_elements[x_1327];
bin_alloc = x_1329;
let x_1332 : u32 = bin_alloc.offset;
inst_ref = BinInstanceRef(x_1332);
let x_1336 : BinInstanceRef = inst_ref;
param_179 = x_1336;
let x_1338 : u32 = ix_11;
param_180 = x_1338;
let x_1339 : BinInstanceRef = BinInstance_index_struct_BinInstanceRef_u11_u1_(&(param_179), &(param_180));
let x_1341 : Alloc = bin_alloc;
param_181 = x_1341;
param_182 = x_1339;
let x_1343 : BinInstance = BinInstance_read_struct_Alloc_u11_struct_BinInstanceRef_u11_(&(param_181), &(param_182));
inst = x_1343;
let x_1345 : u32 = th_ix;
let x_1347 : u32 = inst.element_ix;
sh_elements[x_1345] = x_1347;
}
workgroupBarrier();
let x_1349 : u32 = rd_ix;
let x_1351 : u32 = ready_ix;
wr_ix = min((x_1349 + 256u), x_1351);
continuing {
var x_1365 : bool;
var x_1366 : bool;
let x_1353 : u32 = wr_ix;
let x_1354 : u32 = rd_ix;
let x_1356 : bool = ((x_1353 - x_1354) < 256u);
x_1366 = x_1356;
if (x_1356) {
let x_1359 : u32 = wr_ix;
let x_1360 : u32 = ready_ix;
let x_1362 : u32 = partition_ix;
let x_1363 : u32 = n_partitions;
x_1365 = ((x_1359 < x_1360) | (x_1362 < x_1363));
x_1366 = x_1365;
}
if (x_1366) {
} else {
break;
}
}
}
tag = 0u;
let x_1368 : u32 = th_ix;
let x_1369 : u32 = rd_ix;
let x_1371 : u32 = wr_ix;
if (((x_1368 + x_1369) < x_1371)) {
let x_1376 : u32 = th_ix;
let x_1378 : u32 = sh_elements[x_1376];
element_ix = x_1378;
let x_1383 : u32 = drawtag_start;
let x_1384 : u32 = element_ix;
let x_1387 : u32 = x_1382.scene[(x_1383 + x_1384)];
tag = x_1387;
}
let x_1388 : u32 = tag;
switch(x_1388) {
case 5u, 37u, 68u, 72u, 276u, 732u: {
let x_1393 : u32 = drawmonoid_start;
let x_1394 : u32 = element_ix;
drawmonoid_base = (x_1393 + (4u * x_1394));
let x_1398 : u32 = drawmonoid_base;
let x_1400 : u32 = x_270.memory[x_1398];
path_ix = x_1400;
let x_1403 : u32 = x_883.conf.tile_alloc.offset;
let x_1404 : u32 = path_ix;
let x_1410 : Alloc_1 = x_883.conf.tile_alloc;
param_183.offset = x_1410.offset;
param_184 = PathRef((x_1403 + (x_1404 * 12u)));
let x_1414 : Path = Path_read_struct_Alloc_u11_struct_PathRef_u11_(&(param_183), &(param_184));
path = x_1414;
let x_1417 : u32 = path.bbox.z;
let x_1419 : u32 = path.bbox.x;
stride = (x_1417 - x_1419);
let x_1422 : u32 = th_ix;
let x_1423 : u32 = stride;
sh_tile_stride[x_1422] = x_1423;
let x_1427 : u32 = path.bbox.x;
let x_1429 : u32 = bin_tile_x;
dx = (bitcast<i32>(x_1427) - bitcast<i32>(x_1429));
let x_1434 : u32 = path.bbox.y;
let x_1436 : u32 = bin_tile_y;
dy = (bitcast<i32>(x_1434) - bitcast<i32>(x_1436));
let x_1440 : i32 = dx;
x0 = clamp(x_1440, 0i, 16i);
let x_1443 : i32 = dy;
y0 = clamp(x_1443, 0i, 16i);
let x_1447 : u32 = path.bbox.z;
let x_1449 : u32 = bin_tile_x;
x1 = clamp((bitcast<i32>(x_1447) - bitcast<i32>(x_1449)), 0i, 16i);
let x_1455 : u32 = path.bbox.w;
let x_1457 : u32 = bin_tile_y;
y1 = clamp((bitcast<i32>(x_1455) - bitcast<i32>(x_1457)), 0i, 16i);
let x_1462 : u32 = th_ix;
let x_1463 : i32 = x1;
let x_1464 : i32 = x0;
sh_tile_width[x_1462] = bitcast<u32>((x_1463 - x_1464));
let x_1469 : u32 = th_ix;
let x_1470 : i32 = x0;
sh_tile_x0[x_1469] = bitcast<u32>(x_1470);
let x_1474 : u32 = th_ix;
let x_1475 : i32 = y0;
sh_tile_y0[x_1474] = bitcast<u32>(x_1475);
let x_1479 : i32 = x1;
let x_1480 : i32 = x0;
let x_1483 : i32 = y1;
let x_1484 : i32 = y0;
tile_count = (bitcast<u32>((x_1479 - x_1480)) * bitcast<u32>((x_1483 - x_1484)));
let x_1490 : u32 = path.tiles.offset;
let x_1491 : i32 = dy;
let x_1493 : u32 = stride;
let x_1495 : i32 = dx;
base = (x_1490 - (((bitcast<u32>(x_1491) * x_1493) + bitcast<u32>(x_1495)) * 8u));
let x_1501 : u32 = th_ix;
let x_1502 : u32 = base;
sh_tile_base[x_1501] = x_1502;
let x_1506 : u32 = path.bbox.z;
let x_1508 : u32 = path.bbox.x;
let x_1511 : u32 = path.bbox.w;
let x_1513 : u32 = path.bbox.y;
let x_1519 : u32 = path.tiles.offset;
param_185 = x_1519;
param_186 = (((x_1506 - x_1508) * (x_1511 - x_1513)) * 8u);
param_187 = true;
let x_1522 : Alloc = new_alloc_u1_u1_b1_(&(param_185), &(param_186), &(param_187));
path_alloc = x_1522;
let x_1524 : u32 = th_ix;
param_188 = x_1524;
let x_1526 : Alloc = path_alloc;
param_189 = x_1526;
write_tile_alloc_u1_struct_Alloc_u11_(&(param_188), &(param_189));
}
default: {
tile_count = 0u;
}
}
let x_1532 : u32 = th_ix;
let x_1533 : u32 = tile_count;
sh_tile_count[x_1532] = x_1533;
i_3 = 0u;
loop {
let x_1541 : u32 = i_3;
if ((x_1541 < 8u)) {
} else {
break;
}
workgroupBarrier();
let x_1543 : u32 = th_ix;
let x_1544 : u32 = i_3;
if ((x_1543 >= (1u << x_1544))) {
let x_1549 : u32 = th_ix;
let x_1550 : u32 = i_3;
let x_1554 : u32 = sh_tile_count[(x_1549 - (1u << x_1550))];
let x_1555 : u32 = tile_count;
tile_count = (x_1555 + x_1554);
}
workgroupBarrier();
let x_1557 : u32 = th_ix;
let x_1558 : u32 = tile_count;
sh_tile_count[x_1557] = x_1558;
continuing {
let x_1560 : u32 = i_3;
i_3 = (x_1560 + bitcast<u32>(1i));
}
}
workgroupBarrier();
let x_1564 : u32 = sh_tile_count[255i];
total_tile_count = x_1564;
let x_1566 : u32 = th_ix;
ix_12 = x_1566;
loop {
var x_1706 : bool;
var x_1707 : bool;
let x_1572 : u32 = ix_12;
let x_1573 : u32 = total_tile_count;
if ((x_1572 < x_1573)) {
} else {
break;
}
el_ix_2 = 0u;
i_4 = 0u;
loop {
let x_1582 : u32 = i_4;
if ((x_1582 < 8u)) {
} else {
break;
}
let x_1585 : u32 = el_ix_2;
let x_1586 : u32 = i_4;
probe_1 = (x_1585 + (128u >> x_1586));
let x_1589 : u32 = ix_12;
let x_1590 : u32 = probe_1;
let x_1593 : u32 = sh_tile_count[(x_1590 - 1u)];
if ((x_1589 >= x_1593)) {
let x_1597 : u32 = probe_1;
el_ix_2 = x_1597;
}
continuing {
let x_1598 : u32 = i_4;
i_4 = (x_1598 + bitcast<u32>(1i));
}
}
let x_1601 : u32 = el_ix_2;
let x_1603 : u32 = sh_elements[x_1601];
element_ix_1 = x_1603;
let x_1605 : u32 = drawtag_start;
let x_1606 : u32 = element_ix_1;
let x_1609 : u32 = x_1382.scene[(x_1605 + x_1606)];
tag_1 = x_1609;
let x_1611 : u32 = ix_12;
let x_1612 : u32 = el_ix_2;
if ((x_1612 > 0u)) {
let x_1617 : u32 = el_ix_2;
let x_1620 : u32 = sh_tile_count[(x_1617 - 1u)];
x_1614 = x_1620;
} else {
x_1614 = 0u;
}
let x_1622 : u32 = x_1614;
seq_ix = (x_1611 - x_1622);
let x_1625 : u32 = el_ix_2;
let x_1627 : u32 = sh_tile_width[x_1625];
width = x_1627;
let x_1629 : u32 = el_ix_2;
let x_1631 : u32 = sh_tile_x0[x_1629];
let x_1632 : u32 = seq_ix;
let x_1633 : u32 = width;
x = (x_1631 + (x_1632 % x_1633));
let x_1637 : u32 = el_ix_2;
let x_1639 : u32 = sh_tile_y0[x_1637];
let x_1640 : u32 = seq_ix;
let x_1641 : u32 = width;
y = (x_1639 + (x_1640 / x_1641));
include_tile = false;
let x_1647 : u32 = el_ix_2;
param_190 = x_1647;
param_191 = true;
let x_1649 : Alloc = read_tile_alloc_u1_b1_(&(param_190), &(param_191));
let x_1650 : u32 = el_ix_2;
let x_1652 : u32 = sh_tile_base[x_1650];
let x_1653 : u32 = el_ix_2;
let x_1655 : u32 = sh_tile_stride[x_1653];
let x_1656 : u32 = y;
let x_1658 : u32 = x;
param_192 = x_1649;
param_193 = TileRef((x_1652 + (((x_1655 * x_1656) + x_1658) * 8u)));
let x_1665 : Tile = Tile_read_struct_Alloc_u11_struct_TileRef_u11_(&(param_192), &(param_193));
tile_1 = x_1665;
let x_1667 : u32 = tag_1;
is_clip = ((x_1667 & 1u) != 0u);
is_blend = false;
let x_1671 : bool = is_clip;
if (x_1671) {
let x_1675 : u32 = drawmonoid_start;
let x_1676 : u32 = element_ix_1;
drawmonoid_base_1 = (x_1675 + (4u * x_1676));
let x_1680 : u32 = drawmonoid_base_1;
let x_1683 : u32 = x_270.memory[(x_1680 + 2u)];
scene_offset = x_1683;
let x_1685 : u32 = drawdata_start;
let x_1686 : u32 = scene_offset;
dd = (x_1685 + (x_1686 >> bitcast<u32>(2i)));
let x_1690 : u32 = dd;
let x_1692 : u32 = x_1382.scene[x_1690];
blend = x_1692;
let x_1693 : u32 = blend;
is_blend = (x_1693 != 32771u);
}
let x_1697 : u32 = tile_1.tile.offset;
let x_1698 : bool = (x_1697 != 0u);
x_1707 = x_1698;
if (!(x_1698)) {
let x_1703 : i32 = tile_1.backdrop;
let x_1705 : bool = is_clip;
x_1706 = ((x_1703 == 0i) == x_1705);
x_1707 = x_1706;
}
let x_1708 : bool = is_blend;
include_tile = (x_1707 | x_1708);
let x_1710 : bool = include_tile;
if (x_1710) {
let x_1714 : u32 = el_ix_2;
el_slice = (x_1714 / 32u);
let x_1718 : u32 = el_ix_2;
el_mask = (1u << (x_1718 & 31u));
}
continuing {
let x_1722 : u32 = ix_12;
ix_12 = (x_1722 + 256u);
}
}
workgroupBarrier();
slice_ix = 0u;
let x_1726 : u32 = th_ix;
let x_1728 : u32 = sh_bitmaps[0i][x_1726];
bitmap = x_1728;
loop {
if (true) {
} else {
break;
}
let x_1734 : u32 = bitmap;
if ((x_1734 == 0u)) {
let x_1738 : u32 = slice_ix;
slice_ix = (x_1738 + bitcast<u32>(1i));
let x_1740 : u32 = slice_ix;
if ((x_1740 == 8u)) {
break;
}
let x_1745 : u32 = slice_ix;
let x_1746 : u32 = th_ix;
let x_1748 : u32 = sh_bitmaps[x_1745][x_1746];
bitmap = x_1748;
let x_1749 : u32 = bitmap;
if ((x_1749 == 0u)) {
continue;
}
}
let x_1755 : u32 = slice_ix;
element_ref_ix = (x_1755 * 32u) + firstTrailingBit(bitmap);
let x_1758 : u32 = element_ref_ix;
let x_1760 : u32 = sh_elements[x_1758];
element_ix_2 = x_1760;
let x_1761 : u32 = bitmap;
let x_1763 : u32 = bitmap;
bitmap = (x_1763 & (x_1761 - 1u));
let x_1766 : u32 = drawtag_start;
let x_1767 : u32 = element_ix_2;
let x_1770 : u32 = x_1382.scene[(x_1766 + x_1767)];
drawtag = x_1770;
let x_1771 : u32 = clip_zero_depth;
if ((x_1771 == 0u)) {
var x_2105 : bool;
var x_2106 : bool;
let x_1777 : u32 = element_ref_ix;
param_194 = x_1777;
param_195 = true;
let x_1779 : Alloc = read_tile_alloc_u1_b1_(&(param_194), &(param_195));
let x_1780 : u32 = element_ref_ix;
let x_1782 : u32 = sh_tile_base[x_1780];
let x_1783 : u32 = element_ref_ix;
let x_1785 : u32 = sh_tile_stride[x_1783];
let x_1786 : u32 = tile_y;
let x_1788 : u32 = tile_x;
param_196 = x_1779;
param_197 = TileRef((x_1782 + (((x_1785 * x_1786) + x_1788) * 8u)));
let x_1795 : Tile = Tile_read_struct_Alloc_u11_struct_TileRef_u11_(&(param_196), &(param_197));
tile_2 = x_1795;
let x_1797 : u32 = drawmonoid_start;
let x_1798 : u32 = element_ix_2;
drawmonoid_base_2 = (x_1797 + (4u * x_1798));
let x_1802 : u32 = drawmonoid_base_2;
let x_1805 : u32 = x_270.memory[(x_1802 + 2u)];
scene_offset_1 = x_1805;
let x_1807 : u32 = drawmonoid_base_2;
let x_1810 : u32 = x_270.memory[(x_1807 + 3u)];
info_offset = x_1810;
let x_1812 : u32 = drawdata_start;
let x_1813 : u32 = scene_offset_1;
dd_1 = (x_1812 + (x_1813 >> bitcast<u32>(2i)));
let x_1817 : u32 = drawinfo_start;
let x_1818 : u32 = info_offset;
di = (x_1817 + (x_1818 >> bitcast<u32>(2i)));
let x_1821 : u32 = drawtag;
switch(x_1821) {
case 37u: {
let x_2142 : u32 = clip_depth;
clip_depth = (x_2142 - bitcast<u32>(1i));
let x_2146 : Alloc = cmd_alloc_1;
param_243 = x_2146;
let x_2148 : CmdRef = cmd_ref_2;
param_244 = x_2148;
let x_2150 : Tile = tile_2;
param_245 = x_2150;
param_246 = -1.0f;
write_fill_struct_Alloc_u11_struct_CmdRef_u11_struct_Tile_struct_TileSegRef_u11_i11_f1_(&(param_243), &(param_244), &(param_245), &(param_246));
let x_2153 : CmdRef = param_244;
cmd_ref_2 = x_2153;
let x_2155 : u32 = dd_1;
let x_2157 : u32 = x_1382.scene[x_2155];
blend_1 = x_2157;
let x_2158 : bool = mem_ok_2;
if (x_2158) {
let x_2161 : u32 = blend_1;
let x_2164 : Alloc = cmd_alloc_1;
param_247 = x_2164;
let x_2166 : CmdRef = cmd_ref_2;
param_248 = x_2166;
param_249 = CmdEndClip(x_2161);
Cmd_EndClip_write_struct_Alloc_u11_struct_CmdRef_u11_struct_CmdEndClip_u11_(&(param_247), &(param_248), &(param_249));
}
let x_2170 : u32 = cmd_ref_2.offset;
cmd_ref_2.offset = (x_2170 + 8u);
let x_2173 : u32 = render_blend_depth;
render_blend_depth = (x_2173 - bitcast<u32>(1i));
}
case 5u: {
let x_2099 : u32 = tile_2.tile.offset;
let x_2100 : bool = (x_2099 == 0u);
x_2106 = x_2100;
if (x_2100) {
let x_2104 : i32 = tile_2.backdrop;
x_2105 = (x_2104 == 0i);
x_2106 = x_2105;
}
if (x_2106) {
let x_2109 : u32 = clip_depth;
clip_zero_depth = (x_2109 + 1u);
} else {
let x_2113 : Alloc = cmd_alloc_1;
param_238 = x_2113;
let x_2115 : CmdRef = cmd_ref_2;
param_239 = x_2115;
let x_2117 : u32 = cmd_limit_1;
param_240 = x_2117;
alloc_cmd_struct_Alloc_u11_struct_CmdRef_u11_u1_(&(param_238), &(param_239), &(param_240));
let x_2119 : Alloc = param_238;
cmd_alloc_1 = x_2119;
let x_2120 : CmdRef = param_239;
cmd_ref_2 = x_2120;
let x_2121 : u32 = param_240;
cmd_limit_1 = x_2121;
let x_2122 : bool = mem_ok_2;
if (x_2122) {
let x_2126 : Alloc = cmd_alloc_1;
param_241 = x_2126;
let x_2128 : CmdRef = cmd_ref_2;
param_242 = x_2128;
Cmd_BeginClip_write_struct_Alloc_u11_struct_CmdRef_u11_(&(param_241), &(param_242));
}
let x_2131 : u32 = cmd_ref_2.offset;
cmd_ref_2.offset = (x_2131 + 4u);
let x_2134 : u32 = render_blend_depth;
render_blend_depth = (x_2134 + bitcast<u32>(1i));
let x_2136 : u32 = max_blend_depth;
let x_2137 : u32 = render_blend_depth;
max_blend_depth = max(x_2136, x_2137);
}
let x_2139 : u32 = clip_depth;
clip_depth = (x_2139 + bitcast<u32>(1i));
}
case 72u: {
let x_2039 : Alloc = cmd_alloc_1;
param_228 = x_2039;
let x_2041 : CmdRef = cmd_ref_2;
param_229 = x_2041;
let x_2043 : u32 = cmd_limit_1;
param_230 = x_2043;
alloc_cmd_struct_Alloc_u11_struct_CmdRef_u11_u1_(&(param_228), &(param_229), &(param_230));
let x_2045 : Alloc = param_228;
cmd_alloc_1 = x_2045;
let x_2046 : CmdRef = param_229;
cmd_ref_2 = x_2046;
let x_2047 : u32 = param_230;
cmd_limit_1 = x_2047;
let x_2048 : u32 = di;
let x_2050 : u32 = x_270.memory[x_2048];
linewidth_1 = bitcast<f32>(x_2050);
let x_2053 : Alloc = cmd_alloc_1;
param_231 = x_2053;
let x_2055 : CmdRef = cmd_ref_2;
param_232 = x_2055;
let x_2057 : Tile = tile_2;
param_233 = x_2057;
let x_2059 : f32 = linewidth_1;
param_234 = x_2059;
write_fill_struct_Alloc_u11_struct_CmdRef_u11_struct_Tile_struct_TileSegRef_u11_i11_f1_(&(param_231), &(param_232), &(param_233), &(param_234));
let x_2061 : CmdRef = param_232;
cmd_ref_2 = x_2061;
let x_2063 : u32 = dd_1;
let x_2065 : u32 = x_1382.scene[x_2063];
index_1 = x_2065;
let x_2067 : u32 = dd_1;
let x_2070 : u32 = x_1382.scene[(x_2067 + 1u)];
raw1_2 = x_2070;
let x_2073 : u32 = raw1_2;
let x_2077 : u32 = raw1_2;
offset_8 = vec2<i32>((bitcast<i32>((x_2073 << bitcast<u32>(16i))) >> bitcast<u32>(16i)), (bitcast<i32>(x_2077) >> bitcast<u32>(16i)));
let x_2081 : bool = mem_ok_2;
if (x_2081) {
let x_2084 : u32 = index_1;
let x_2085 : vec2<i32> = offset_8;
let x_2088 : Alloc = cmd_alloc_1;
param_235 = x_2088;
let x_2090 : CmdRef = cmd_ref_2;
param_236 = x_2090;
param_237 = CmdImage(x_2084, x_2085);
Cmd_Image_write_struct_Alloc_u11_struct_CmdRef_u11_struct_CmdImage_u1_vi21_(&(param_235), &(param_236), &(param_237));
}
let x_2094 : u32 = cmd_ref_2.offset;
cmd_ref_2.offset = (x_2094 + 12u);
}
case 732u: {
let x_1938 : Alloc = cmd_alloc_1;
param_218 = x_1938;
let x_1940 : CmdRef = cmd_ref_2;
param_219 = x_1940;
let x_1942 : u32 = cmd_limit_1;
param_220 = x_1942;
alloc_cmd_struct_Alloc_u11_struct_CmdRef_u11_u1_(&(param_218), &(param_219), &(param_220));
let x_1944 : Alloc = param_218;
cmd_alloc_1 = x_1944;
let x_1945 : CmdRef = param_219;
cmd_ref_2 = x_1945;
let x_1946 : u32 = param_220;
cmd_limit_1 = x_1946;
let x_1947 : u32 = di;
let x_1949 : u32 = x_270.memory[x_1947];
linewidth_1 = bitcast<f32>(x_1949);
let x_1952 : Alloc = cmd_alloc_1;
param_221 = x_1952;
let x_1954 : CmdRef = cmd_ref_2;
param_222 = x_1954;
let x_1956 : Tile = tile_2;
param_223 = x_1956;
let x_1958 : f32 = linewidth_1;
param_224 = x_1958;
write_fill_struct_Alloc_u11_struct_CmdRef_u11_struct_Tile_struct_TileSegRef_u11_i11_f1_(&(param_221), &(param_222), &(param_223), &(param_224));
let x_1960 : CmdRef = param_222;
cmd_ref_2 = x_1960;
let x_1962 : u32 = dd_1;
let x_1964 : u32 = x_1382.scene[x_1962];
cmd_rad.index = x_1964;
let x_1966 : u32 = di;
let x_1969 : u32 = x_270.memory[(x_1966 + 1u)];
let x_1970 : u32 = di;
let x_1973 : u32 = x_270.memory[(x_1970 + 2u)];
let x_1974 : u32 = di;
let x_1977 : u32 = x_270.memory[(x_1974 + 3u)];
let x_1978 : u32 = di;
let x_1981 : u32 = x_270.memory[(x_1978 + 4u)];
cmd_rad.mat_ = bitcast<vec4<f32>>(vec4<u32>(x_1969, x_1973, x_1977, x_1981));
let x_1986 : u32 = di;
let x_1989 : u32 = x_270.memory[(x_1986 + 5u)];
let x_1990 : u32 = di;
let x_1993 : u32 = x_270.memory[(x_1990 + 6u)];
cmd_rad.xlat = bitcast<vec2<f32>>(vec2<u32>(x_1989, x_1993));
let x_1999 : u32 = di;
let x_2002 : u32 = x_270.memory[(x_1999 + 7u)];
let x_2003 : u32 = di;
let x_2006 : u32 = x_270.memory[(x_2003 + 8u)];
cmd_rad.c1 = bitcast<vec2<f32>>(vec2<u32>(x_2002, x_2006));
let x_2010 : u32 = di;
let x_2013 : u32 = x_270.memory[(x_2010 + 9u)];
cmd_rad.ra = bitcast<f32>(x_2013);
let x_2016 : u32 = di;
let x_2019 : u32 = x_270.memory[(x_2016 + 10u)];
cmd_rad.roff = bitcast<f32>(x_2019);
let x_2022 : bool = mem_ok_2;
if (x_2022) {
let x_2026 : Alloc = cmd_alloc_1;
param_225 = x_2026;
let x_2028 : CmdRef = cmd_ref_2;
param_226 = x_2028;
let x_2030 : CmdRadGrad = cmd_rad;
param_227 = x_2030;
Cmd_RadGrad_write_struct_Alloc_u11_struct_CmdRef_u11_struct_CmdRadGrad_u1_vf4_vf2_vf2_f1_f11_(&(param_225), &(param_226), &(param_227));
}
let x_2034 : u32 = cmd_ref_2.offset;
cmd_ref_2.offset = (x_2034 + 48u);
}
case 276u: {
let x_1875 : Alloc = cmd_alloc_1;
param_208 = x_1875;
let x_1877 : CmdRef = cmd_ref_2;
param_209 = x_1877;
let x_1879 : u32 = cmd_limit_1;
param_210 = x_1879;
alloc_cmd_struct_Alloc_u11_struct_CmdRef_u11_u1_(&(param_208), &(param_209), &(param_210));
let x_1881 : Alloc = param_208;
cmd_alloc_1 = x_1881;
let x_1882 : CmdRef = param_209;
cmd_ref_2 = x_1882;
let x_1883 : u32 = param_210;
cmd_limit_1 = x_1883;
let x_1884 : u32 = di;
let x_1886 : u32 = x_270.memory[x_1884];
linewidth_1 = bitcast<f32>(x_1886);
let x_1889 : Alloc = cmd_alloc_1;
param_211 = x_1889;
let x_1891 : CmdRef = cmd_ref_2;
param_212 = x_1891;
let x_1893 : Tile = tile_2;
param_213 = x_1893;
let x_1895 : f32 = linewidth_1;
param_214 = x_1895;
write_fill_struct_Alloc_u11_struct_CmdRef_u11_struct_Tile_struct_TileSegRef_u11_i11_f1_(&(param_211), &(param_212), &(param_213), &(param_214));
let x_1897 : CmdRef = param_212;
cmd_ref_2 = x_1897;
let x_1899 : u32 = dd_1;
let x_1901 : u32 = x_1382.scene[x_1899];
cmd_lin.index = x_1901;
let x_1903 : u32 = di;
let x_1906 : u32 = x_270.memory[(x_1903 + 1u)];
cmd_lin.line_x = bitcast<f32>(x_1906);
let x_1909 : u32 = di;
let x_1912 : u32 = x_270.memory[(x_1909 + 2u)];
cmd_lin.line_y = bitcast<f32>(x_1912);
let x_1915 : u32 = di;
let x_1918 : u32 = x_270.memory[(x_1915 + 3u)];
cmd_lin.line_c = bitcast<f32>(x_1918);
let x_1921 : bool = mem_ok_2;
if (x_1921) {
let x_1925 : Alloc = cmd_alloc_1;
param_215 = x_1925;
let x_1927 : CmdRef = cmd_ref_2;
param_216 = x_1927;
let x_1929 : CmdLinGrad = cmd_lin;
param_217 = x_1929;
Cmd_LinGrad_write_struct_Alloc_u11_struct_CmdRef_u11_struct_CmdLinGrad_u1_f1_f1_f11_(&(param_215), &(param_216), &(param_217));
}
let x_1933 : u32 = cmd_ref_2.offset;
cmd_ref_2.offset = (x_1933 + 20u);
}
case 68u: {
let x_1830 : u32 = di;
let x_1832 : u32 = x_270.memory[x_1830];
linewidth_1 = bitcast<f32>(x_1832);
let x_1835 : Alloc = cmd_alloc_1;
param_198 = x_1835;
let x_1837 : CmdRef = cmd_ref_2;
param_199 = x_1837;
let x_1839 : u32 = cmd_limit_1;
param_200 = x_1839;
alloc_cmd_struct_Alloc_u11_struct_CmdRef_u11_u1_(&(param_198), &(param_199), &(param_200));
let x_1841 : Alloc = param_198;
cmd_alloc_1 = x_1841;
let x_1842 : CmdRef = param_199;
cmd_ref_2 = x_1842;
let x_1843 : u32 = param_200;
cmd_limit_1 = x_1843;
let x_1845 : Alloc = cmd_alloc_1;
param_201 = x_1845;
let x_1847 : CmdRef = cmd_ref_2;
param_202 = x_1847;
let x_1849 : Tile = tile_2;
param_203 = x_1849;
let x_1851 : f32 = linewidth_1;
param_204 = x_1851;
write_fill_struct_Alloc_u11_struct_CmdRef_u11_struct_Tile_struct_TileSegRef_u11_i11_f1_(&(param_201), &(param_202), &(param_203), &(param_204));
let x_1853 : CmdRef = param_202;
cmd_ref_2 = x_1853;
let x_1855 : u32 = dd_1;
let x_1857 : u32 = x_1382.scene[x_1855];
rgba = x_1857;
let x_1858 : bool = mem_ok_2;
if (x_1858) {
let x_1861 : u32 = rgba;
let x_1864 : Alloc = cmd_alloc_1;
param_205 = x_1864;
let x_1866 : CmdRef = cmd_ref_2;
param_206 = x_1866;
param_207 = CmdColor(x_1861);
Cmd_Color_write_struct_Alloc_u11_struct_CmdRef_u11_struct_CmdColor_u11_(&(param_205), &(param_206), &(param_207));
}
let x_1870 : u32 = cmd_ref_2.offset;
cmd_ref_2.offset = (x_1870 + 8u);
}
default: {
}
}
} else {
let x_2178 : u32 = drawtag;
switch(x_2178) {
case 37u: {
let x_2185 : u32 = clip_depth;
let x_2186 : u32 = clip_zero_depth;
if ((x_2185 == x_2186)) {
clip_zero_depth = 0u;
}
let x_2190 : u32 = clip_depth;
clip_depth = (x_2190 - bitcast<u32>(1i));
}
case 5u: {
let x_2182 : u32 = clip_depth;
clip_depth = (x_2182 + bitcast<u32>(1i));
}
default: {
}
}
}
}
workgroupBarrier();
let x_2194 : u32 = rd_ix;
rd_ix = (x_2194 + 256u);
let x_2196 : u32 = rd_ix;
let x_2197 : u32 = ready_ix;
let x_2199 : u32 = partition_ix;
let x_2200 : u32 = n_partitions;
if (((x_2196 >= x_2197) & (x_2199 >= x_2200))) {
break;
}
}
let x_2206 : u32 = bin_tile_x;
let x_2207 : u32 = tile_x;
let x_2210 : u32 = x_883.conf.width_in_tiles;
let x_2211 : bool = ((x_2206 + x_2207) < x_2210);
x_2220 = x_2211;
if (x_2211) {
let x_2214 : u32 = bin_tile_y;
let x_2215 : u32 = tile_y;
let x_2218 : u32 = x_883.conf.height_in_tiles;
x_2219 = ((x_2214 + x_2215) < x_2218);
x_2220 = x_2219;
}
if (x_2220) {
let x_2223 : bool = mem_ok_2;
if (x_2223) {
let x_2227 : Alloc = cmd_alloc_1;
param_250 = x_2227;
let x_2229 : CmdRef = cmd_ref_2;
param_251 = x_2229;
Cmd_End_write_struct_Alloc_u11_struct_CmdRef_u11_(&(param_250), &(param_251));
}
let x_2231 : u32 = max_blend_depth;
if ((x_2231 > 4u)) {
let x_2236 : u32 = max_blend_depth;
scratch_size = ((((x_2236 * 16u) * 16u) * 1u) * 4u);
let x_2243 : u32 = scratch_size;
let x_2244 : u32 = 0u;
scratch = x_2244;
let x_2246 : u32 = scratch_alloc.offset;
let x_2249 : Alloc = scratch_alloc;
param_252 = x_2249;
param_253 = (x_2246 >> bitcast<u32>(2i));
let x_2252 : u32 = scratch;
param_254 = x_2252;
write_mem_struct_Alloc_u11_u1_u1_(&(param_252), &(param_253), &(param_254));
}
}
return;
}
@compute @workgroup_size(256i, 1i, 1i)
fn main(@builtin(workgroup_id) gl_WorkGroupID_param : vec3<u32>, @builtin(local_invocation_id) gl_LocalInvocationID_param : vec3<u32>) {
gl_WorkGroupID = gl_WorkGroupID_param;
gl_LocalInvocationID = gl_LocalInvocationID_param;
main_1();
}
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment