Created
October 7, 2022 01:41
-
-
Save raphlinus/1406c2a41589ea1a616cfc2552d6dc1c to your computer and use it in GitHub Desktop.
Partially translated coarse.comp to wgsl
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
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