alex/vello
1
0
Fork 0
mirror of https://github.com/italicsjenga/vello.git synced 2025-01-10 20:51:29 +11:00
Code Issues Packages Projects Releases Wiki Activity

Regenerate generated shaders

Browse source 
This just runs ninja on the piet-gpu/shaders on a Windows machine, so
translated shaders match the existing pipeline.

At some point, we'll rework this to reduce friction.
This commit is contained in:
Raph Levien 2022-03-07 12:49:59 -08:00
parent d3b08e4c52
commit 90774f1f46
15 changed files with 1607 additions and 535 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

10
piet-gpu/shader/gen/binning.msl generated
View file

@ -220,7 +220,7 @@ void BinInstance_write(thread const Alloc& a, thread const BinInstanceRef& ref,
kernel void main0(constant uint* spvBufferSizeConstants [[buffer(25)]], device Memory& v_94 [[buffer(0)]], const device ConfigBuf& v_202 [[buffer(1)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]]) kernel void main0(constant uint* spvBufferSizeConstants [[buffer(25)]], device Memory& v_94 [[buffer(0)]], const device ConfigBuf& v_202 [[buffer(1)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
{ {
threadgroup uint bitmaps[8][256]; threadgroup uint bitmaps[8][256];
threadgroup bool sh_alloc_failed; threadgroup short sh_alloc_failed;
threadgroup uint count[8][256]; threadgroup uint count[8][256];
threadgroup Alloc sh_chunk_alloc[256]; threadgroup Alloc sh_chunk_alloc[256];
constant uint& v_94BufferSize = spvBufferSizeConstants[0]; constant uint& v_94BufferSize = spvBufferSizeConstants[0];
@ -232,7 +232,7 @@ kernel void main0(constant uint* spvBufferSizeConstants [[buffer(25)]], device M
} }
if (gl_LocalInvocationID.x == 0u) if (gl_LocalInvocationID.x == 0u)
{ {
sh_alloc_failed = false; sh_alloc_failed = short(false);
} }
threadgroup_barrier(mem_flags::mem_threadgroup); threadgroup_barrier(mem_flags::mem_threadgroup);
uint element_ix = (my_partition * 256u) + gl_LocalInvocationID.x; uint element_ix = (my_partition * 256u) + gl_LocalInvocationID.x;
@ -331,7 +331,7 @@ kernel void main0(constant uint* spvBufferSizeConstants [[buffer(25)]], device M
sh_chunk_alloc[gl_LocalInvocationID.x] = chunk_alloc; sh_chunk_alloc[gl_LocalInvocationID.x] = chunk_alloc;
if (chunk.failed) if (chunk.failed)
{ {
sh_alloc_failed = true; sh_alloc_failed = short(true);
} }
} }
uint out_ix = (v_202.conf.bin_alloc.offset >> uint(2)) + (((my_partition * 256u) + gl_LocalInvocationID.x) * 2u); uint out_ix = (v_202.conf.bin_alloc.offset >> uint(2)) + (((my_partition * 256u) + gl_LocalInvocationID.x) * 2u);
@ -347,13 +347,13 @@ kernel void main0(constant uint* spvBufferSizeConstants [[buffer(25)]], device M
write_mem(param_16, param_17, param_18, v_94, v_94BufferSize); write_mem(param_16, param_17, param_18, v_94, v_94BufferSize);
threadgroup_barrier(mem_flags::mem_threadgroup); threadgroup_barrier(mem_flags::mem_threadgroup);
bool _687; bool _687;
if (!sh_alloc_failed) if (!bool(sh_alloc_failed))
{ {
_687 = v_94.mem_error != 0u; _687 = v_94.mem_error != 0u;
} }
else else
{ {
_687 = sh_alloc_failed; _687 = bool(sh_alloc_failed);
} }
if (_687) if (_687)
{ {

BIN
piet-gpu/shader/gen/coarse.dxil generated
View file

Binary file not shown.

478
piet-gpu/shader/gen/coarse.hlsl generated
View file

@ -49,6 +49,17 @@ struct AnnoLinGradient
float line_c; float line_c;
}; };
struct AnnoEndClipRef
{
uint offset;
};
struct AnnoEndClip
{
float4 bbox;
uint blend;
};
struct AnnotatedRef struct AnnotatedRef
{ {
uint offset; uint offset;
@ -153,6 +164,16 @@ struct CmdImage
int2 offset; int2 offset;
}; };
struct CmdEndClipRef
{
uint offset;
};
struct CmdEndClip
{
uint blend;
};
struct CmdJumpRef struct CmdJumpRef
{ {
uint offset; uint offset;
@ -197,8 +218,8 @@ struct Config
static const uint3 gl_WorkGroupSize = uint3(256u, 1u, 1u); static const uint3 gl_WorkGroupSize = uint3(256u, 1u, 1u);
RWByteAddressBuffer _283 : register(u0, space0); RWByteAddressBuffer _308 : register(u0, space0);
ByteAddressBuffer _1169 : register(t1, space0); ByteAddressBuffer _1283 : register(t1, space0);
static uint3 gl_WorkGroupID; static uint3 gl_WorkGroupID;
static uint3 gl_LocalInvocationID; static uint3 gl_LocalInvocationID;
@ -221,8 +242,8 @@ groupshared uint sh_tile_count[256];
Alloc slice_mem(Alloc a, uint offset, uint size) Alloc slice_mem(Alloc a, uint offset, uint size)
{ {
Alloc _360 = { a.offset + offset }; Alloc _385 = { a.offset + offset };
return _360; return _385;
} }
bool touch_mem(Alloc alloc, uint offset) bool touch_mem(Alloc alloc, uint offset)
@ -238,7 +259,7 @@ uint read_mem(Alloc alloc, uint offset)
{ {
return 0u; return 0u;
} }
uint v = _283.Load(offset * 4 + 8); uint v = _308.Load(offset * 4 + 8);
return v; return v;
} }
@ -251,8 +272,8 @@ Alloc new_alloc(uint offset, uint size, bool mem_ok)
BinInstanceRef BinInstance_index(BinInstanceRef ref, uint index) BinInstanceRef BinInstance_index(BinInstanceRef ref, uint index)
{ {
BinInstanceRef _674 = { ref.offset + (index * 4u) }; BinInstanceRef _765 = { ref.offset + (index * 4u) };
return _674; return _765;
} }
BinInstance BinInstance_read(Alloc a, BinInstanceRef ref) BinInstance BinInstance_read(Alloc a, BinInstanceRef ref)
@ -271,8 +292,8 @@ AnnotatedTag Annotated_tag(Alloc a, AnnotatedRef ref)
Alloc param = a; Alloc param = a;
uint param_1 = ref.offset >> uint(2); uint param_1 = ref.offset >> uint(2);
uint tag_and_flags = read_mem(param, param_1); uint tag_and_flags = read_mem(param, param_1);
AnnotatedTag _636 = { tag_and_flags & 65535u, tag_and_flags >> uint(16) }; AnnotatedTag _717 = { tag_and_flags & 65535u, tag_and_flags >> uint(16) };
return _636; return _717;
} }
Path Path_read(Alloc a, PathRef ref) Path Path_read(Alloc a, PathRef ref)
@ -289,8 +310,8 @@ Path Path_read(Alloc a, PathRef ref)
uint raw2 = read_mem(param_4, param_5); uint raw2 = read_mem(param_4, param_5);
Path s; Path s;
s.bbox = uint4(raw0 & 65535u, raw0 >> uint(16), raw1 & 65535u, raw1 >> uint(16)); s.bbox = uint4(raw0 & 65535u, raw0 >> uint(16), raw1 & 65535u, raw1 >> uint(16));
TileRef _734 = { raw2 }; TileRef _825 = { raw2 };
s.tiles = _734; s.tiles = _825;
return s; return s;
} }
@ -300,11 +321,11 @@ void write_tile_alloc(uint el_ix, Alloc a)
Alloc read_tile_alloc(uint el_ix, bool mem_ok) Alloc read_tile_alloc(uint el_ix, bool mem_ok)
{ {
uint _1055; uint _1169;
_283.GetDimensions(_1055); _308.GetDimensions(_1169);
_1055 = (_1055 - 8) / 4; _1169 = (_1169 - 8) / 4;
uint param = 0u; uint param = 0u;
uint param_1 = uint(int(_1055) * 4); uint param_1 = uint(int(_1169) * 4);
bool param_2 = mem_ok; bool param_2 = mem_ok;
return new_alloc(param, param_1, param_2); return new_alloc(param, param_1, param_2);
} }
@ -318,9 +339,9 @@ Tile Tile_read(Alloc a, TileRef ref)
Alloc param_2 = a; Alloc param_2 = a;
uint param_3 = ix + 1u; uint param_3 = ix + 1u;
uint raw1 = read_mem(param_2, param_3); uint raw1 = read_mem(param_2, param_3);
TileSegRef _759 = { raw0 }; TileSegRef _850 = { raw0 };
Tile s; Tile s;
s.tile = _759; s.tile = _850;
s.backdrop = int(raw1); s.backdrop = int(raw1);
return s; return s;
} }
@ -355,30 +376,30 @@ AnnoColor AnnoColor_read(Alloc a, AnnoColorRef ref)
AnnoColor Annotated_Color_read(Alloc a, AnnotatedRef ref) AnnoColor Annotated_Color_read(Alloc a, AnnotatedRef ref)
{ {
AnnoColorRef _642 = { ref.offset + 4u }; AnnoColorRef _723 = { ref.offset + 4u };
Alloc param = a; Alloc param = a;
AnnoColorRef param_1 = _642; AnnoColorRef param_1 = _723;
return AnnoColor_read(param, param_1); return AnnoColor_read(param, param_1);
} }
MallocResult malloc(uint size) MallocResult malloc(uint size)
{ {
uint _289; uint _314;
_283.InterlockedAdd(0, size, _289); _308.InterlockedAdd(0, size, _314);
uint offset = _289; uint offset = _314;
uint _296; uint _321;
_283.GetDimensions(_296); _308.GetDimensions(_321);
_296 = (_296 - 8) / 4; _321 = (_321 - 8) / 4;
MallocResult r; MallocResult r;
r.failed = (offset + size) > uint(int(_296) * 4); r.failed = (offset + size) > uint(int(_321) * 4);
uint param = offset; uint param = offset;
uint param_1 = size; uint param_1 = size;
bool param_2 = !r.failed; bool param_2 = !r.failed;
r.alloc = new_alloc(param, param_1, param_2); r.alloc = new_alloc(param, param_1, param_2);
if (r.failed) if (r.failed)
{ {
uint _318; uint _343;
_283.InterlockedMax(4, 1u, _318); _308.InterlockedMax(4, 1u, _343);
return r; return r;
} }
return r; return r;
@ -392,7 +413,7 @@ void write_mem(Alloc alloc, uint offset, uint val)
{ {
return; return;
} }
_283.Store(offset * 4 + 8, val); _308.Store(offset * 4 + 8, val);
} }
void CmdJump_write(Alloc a, CmdJumpRef ref, CmdJump s) void CmdJump_write(Alloc a, CmdJumpRef ref, CmdJump s)
@ -410,9 +431,9 @@ void Cmd_Jump_write(Alloc a, CmdRef ref, CmdJump s)
uint param_1 = ref.offset >> uint(2); uint param_1 = ref.offset >> uint(2);
uint param_2 = 10u; uint param_2 = 10u;
write_mem(param, param_1, param_2); write_mem(param, param_1, param_2);
CmdJumpRef _1048 = { ref.offset + 4u }; CmdJumpRef _1162 = { ref.offset + 4u };
Alloc param_3 = a; Alloc param_3 = a;
CmdJumpRef param_4 = _1048; CmdJumpRef param_4 = _1162;
CmdJump param_5 = s; CmdJump param_5 = s;
CmdJump_write(param_3, param_4, param_5); CmdJump_write(param_3, param_4, param_5);
} }
@ -424,21 +445,21 @@ bool alloc_cmd(inout Alloc cmd_alloc, inout CmdRef cmd_ref, inout uint cmd_limit
return true; return true;
} }
uint param = 1024u; uint param = 1024u;
MallocResult _1076 = malloc(param); MallocResult _1190 = malloc(param);
MallocResult new_cmd = _1076; MallocResult new_cmd = _1190;
if (new_cmd.failed) if (new_cmd.failed)
{ {
return false; return false;
} }
CmdJump _1086 = { new_cmd.alloc.offset }; CmdJump _1200 = { new_cmd.alloc.offset };
CmdJump jump = _1086; CmdJump jump = _1200;
Alloc param_1 = cmd_alloc; Alloc param_1 = cmd_alloc;
CmdRef param_2 = cmd_ref; CmdRef param_2 = cmd_ref;
CmdJump param_3 = jump; CmdJump param_3 = jump;
Cmd_Jump_write(param_1, param_2, param_3); Cmd_Jump_write(param_1, param_2, param_3);
cmd_alloc = new_cmd.alloc; cmd_alloc = new_cmd.alloc;
CmdRef _1098 = { cmd_alloc.offset }; CmdRef _1212 = { cmd_alloc.offset };
cmd_ref = _1098; cmd_ref = _1212;
cmd_limit = (cmd_alloc.offset + 1024u) - 60u; cmd_limit = (cmd_alloc.offset + 1024u) - 60u;
return true; return true;
} }
@ -467,9 +488,9 @@ void Cmd_Fill_write(Alloc a, CmdRef ref, CmdFill s)
uint param_1 = ref.offset >> uint(2); uint param_1 = ref.offset >> uint(2);
uint param_2 = 1u; uint param_2 = 1u;
write_mem(param, param_1, param_2); write_mem(param, param_1, param_2);
CmdFillRef _932 = { ref.offset + 4u }; CmdFillRef _1036 = { ref.offset + 4u };
Alloc param_3 = a; Alloc param_3 = a;
CmdFillRef param_4 = _932; CmdFillRef param_4 = _1036;
CmdFill param_5 = s; CmdFill param_5 = s;
CmdFill_write(param_3, param_4, param_5); CmdFill_write(param_3, param_4, param_5);
} }
@ -501,9 +522,9 @@ void Cmd_Stroke_write(Alloc a, CmdRef ref, CmdStroke s)
uint param_1 = ref.offset >> uint(2); uint param_1 = ref.offset >> uint(2);
uint param_2 = 2u; uint param_2 = 2u;
write_mem(param, param_1, param_2); write_mem(param, param_1, param_2);
CmdStrokeRef _950 = { ref.offset + 4u }; CmdStrokeRef _1054 = { ref.offset + 4u };
Alloc param_3 = a; Alloc param_3 = a;
CmdStrokeRef param_4 = _950; CmdStrokeRef param_4 = _1054;
CmdStroke param_5 = s; CmdStroke param_5 = s;
CmdStroke_write(param_3, param_4, param_5); CmdStroke_write(param_3, param_4, param_5);
} }
@ -515,8 +536,8 @@ void write_fill(Alloc alloc, inout CmdRef cmd_ref, uint flags, Tile tile, float
{ {
if (tile.tile.offset != 0u) if (tile.tile.offset != 0u)
{ {
CmdFill _1122 = { tile.tile.offset, tile.backdrop }; CmdFill _1236 = { tile.tile.offset, tile.backdrop };
CmdFill cmd_fill = _1122; CmdFill cmd_fill = _1236;
Alloc param_1 = alloc; Alloc param_1 = alloc;
CmdRef param_2 = cmd_ref; CmdRef param_2 = cmd_ref;
CmdFill param_3 = cmd_fill; CmdFill param_3 = cmd_fill;
@ -533,8 +554,8 @@ void write_fill(Alloc alloc, inout CmdRef cmd_ref, uint flags, Tile tile, float
} }
else else
{ {
CmdStroke _1152 = { tile.tile.offset, 0.5f * linewidth }; CmdStroke _1266 = { tile.tile.offset, 0.5f * linewidth };
CmdStroke cmd_stroke = _1152; CmdStroke cmd_stroke = _1266;
Alloc param_6 = alloc; Alloc param_6 = alloc;
CmdRef param_7 = cmd_ref; CmdRef param_7 = cmd_ref;
CmdStroke param_8 = cmd_stroke; CmdStroke param_8 = cmd_stroke;
@ -558,9 +579,9 @@ void Cmd_Color_write(Alloc a, CmdRef ref, CmdColor s)
uint param_1 = ref.offset >> uint(2); uint param_1 = ref.offset >> uint(2);
uint param_2 = 5u; uint param_2 = 5u;
write_mem(param, param_1, param_2); write_mem(param, param_1, param_2);
CmdColorRef _976 = { ref.offset + 4u }; CmdColorRef _1080 = { ref.offset + 4u };
Alloc param_3 = a; Alloc param_3 = a;
CmdColorRef param_4 = _976; CmdColorRef param_4 = _1080;
CmdColor param_5 = s; CmdColor param_5 = s;
CmdColor_write(param_3, param_4, param_5); CmdColor_write(param_3, param_4, param_5);
} }
@ -607,9 +628,9 @@ AnnoLinGradient AnnoLinGradient_read(Alloc a, AnnoLinGradientRef ref)
AnnoLinGradient Annotated_LinGradient_read(Alloc a, AnnotatedRef ref) AnnoLinGradient Annotated_LinGradient_read(Alloc a, AnnotatedRef ref)
{ {
AnnoLinGradientRef _652 = { ref.offset + 4u }; AnnoLinGradientRef _733 = { ref.offset + 4u };
Alloc param = a; Alloc param = a;
AnnoLinGradientRef param_1 = _652; AnnoLinGradientRef param_1 = _733;
return AnnoLinGradient_read(param, param_1); return AnnoLinGradient_read(param, param_1);
} }
@ -640,9 +661,9 @@ void Cmd_LinGrad_write(Alloc a, CmdRef ref, CmdLinGrad s)
uint param_1 = ref.offset >> uint(2); uint param_1 = ref.offset >> uint(2);
uint param_2 = 6u; uint param_2 = 6u;
write_mem(param, param_1, param_2); write_mem(param, param_1, param_2);
CmdLinGradRef _994 = { ref.offset + 4u }; CmdLinGradRef _1098 = { ref.offset + 4u };
Alloc param_3 = a; Alloc param_3 = a;
CmdLinGradRef param_4 = _994; CmdLinGradRef param_4 = _1098;
CmdLinGrad param_5 = s; CmdLinGrad param_5 = s;
CmdLinGrad_write(param_3, param_4, param_5); CmdLinGrad_write(param_3, param_4, param_5);
} }
@ -681,9 +702,9 @@ AnnoImage AnnoImage_read(Alloc a, AnnoImageRef ref)
AnnoImage Annotated_Image_read(Alloc a, AnnotatedRef ref) AnnoImage Annotated_Image_read(Alloc a, AnnotatedRef ref)
{ {
AnnoImageRef _662 = { ref.offset + 4u }; AnnoImageRef _743 = { ref.offset + 4u };
Alloc param = a; Alloc param = a;
AnnoImageRef param_1 = _662; AnnoImageRef param_1 = _743;
return AnnoImage_read(param, param_1); return AnnoImage_read(param, param_1);
} }
@ -706,9 +727,9 @@ void Cmd_Image_write(Alloc a, CmdRef ref, CmdImage s)
uint param_1 = ref.offset >> uint(2); uint param_1 = ref.offset >> uint(2);
uint param_2 = 7u; uint param_2 = 7u;
write_mem(param, param_1, param_2); write_mem(param, param_1, param_2);
CmdImageRef _1012 = { ref.offset + 4u }; CmdImageRef _1116 = { ref.offset + 4u };
Alloc param_3 = a; Alloc param_3 = a;
CmdImageRef param_4 = _1012; CmdImageRef param_4 = _1116;
CmdImage param_5 = s; CmdImage param_5 = s;
CmdImage_write(param_3, param_4, param_5); CmdImage_write(param_3, param_4, param_5);
} }
@ -721,12 +742,58 @@ void Cmd_BeginClip_write(Alloc a, CmdRef ref)
write_mem(param, param_1, param_2); write_mem(param, param_1, param_2);
} }
void Cmd_EndClip_write(Alloc a, CmdRef ref) AnnoEndClip AnnoEndClip_read(Alloc a, AnnoEndClipRef ref)
{
uint ix = ref.offset >> uint(2);
Alloc param = a;
uint param_1 = ix + 0u;
uint raw0 = read_mem(param, param_1);
Alloc param_2 = a;
uint param_3 = ix + 1u;
uint raw1 = read_mem(param_2, param_3);
Alloc param_4 = a;
uint param_5 = ix + 2u;
uint raw2 = read_mem(param_4, param_5);
Alloc param_6 = a;
uint param_7 = ix + 3u;
uint raw3 = read_mem(param_6, param_7);
Alloc param_8 = a;
uint param_9 = ix + 4u;
uint raw4 = read_mem(param_8, param_9);
AnnoEndClip s;
s.bbox = float4(asfloat(raw0), asfloat(raw1), asfloat(raw2), asfloat(raw3));
s.blend = raw4;
return s;
}
AnnoEndClip Annotated_EndClip_read(Alloc a, AnnotatedRef ref)
{
AnnoEndClipRef _753 = { ref.offset + 4u };
Alloc param = a;
AnnoEndClipRef param_1 = _753;
return AnnoEndClip_read(param, param_1);
}
void CmdEndClip_write(Alloc a, CmdEndClipRef ref, CmdEndClip s)
{
uint ix = ref.offset >> uint(2);
Alloc param = a;
uint param_1 = ix + 0u;
uint param_2 = s.blend;
write_mem(param, param_1, param_2);
}
void Cmd_EndClip_write(Alloc a, CmdRef ref, CmdEndClip s)
{ {
Alloc param = a; Alloc param = a;
uint param_1 = ref.offset >> uint(2); uint param_1 = ref.offset >> uint(2);
uint param_2 = 9u; uint param_2 = 9u;
write_mem(param, param_1, param_2); write_mem(param, param_1, param_2);
CmdEndClipRef _1143 = { ref.offset + 4u };
Alloc param_3 = a;
CmdEndClipRef param_4 = _1143;
CmdEndClip param_5 = s;
CmdEndClip_write(param_3, param_4, param_5);
} }
void Cmd_End_write(Alloc a, CmdRef ref) void Cmd_End_write(Alloc a, CmdRef ref)
@ -739,25 +806,25 @@ void Cmd_End_write(Alloc a, CmdRef ref)
void comp_main() void comp_main()
{ {
uint width_in_bins = ((_1169.Load(8) + 16u) - 1u) / 16u; uint width_in_bins = ((_1283.Load(8) + 16u) - 1u) / 16u;
uint bin_ix = (width_in_bins * gl_WorkGroupID.y) + gl_WorkGroupID.x; uint bin_ix = (width_in_bins * gl_WorkGroupID.y) + gl_WorkGroupID.x;
uint partition_ix = 0u; uint partition_ix = 0u;
uint n_partitions = ((_1169.Load(0) + 256u) - 1u) / 256u; uint n_partitions = ((_1283.Load(0) + 256u) - 1u) / 256u;
uint th_ix = gl_LocalInvocationID.x; uint th_ix = gl_LocalInvocationID.x;
uint bin_tile_x = 16u * gl_WorkGroupID.x; uint bin_tile_x = 16u * gl_WorkGroupID.x;
uint bin_tile_y = 16u * gl_WorkGroupID.y; uint bin_tile_y = 16u * gl_WorkGroupID.y;
uint tile_x = gl_LocalInvocationID.x % 16u; uint tile_x = gl_LocalInvocationID.x % 16u;
uint tile_y = gl_LocalInvocationID.x / 16u; uint tile_y = gl_LocalInvocationID.x / 16u;
uint this_tile_ix = (((bin_tile_y + tile_y) * _1169.Load(8)) + bin_tile_x) + tile_x; uint this_tile_ix = (((bin_tile_y + tile_y) * _1283.Load(8)) + bin_tile_x) + tile_x;
Alloc _1234; Alloc _1348;
_1234.offset = _1169.Load(24); _1348.offset = _1283.Load(24);
Alloc param; Alloc param;
param.offset = _1234.offset; param.offset = _1348.offset;
uint param_1 = this_tile_ix * 1024u; uint param_1 = this_tile_ix * 1024u;
uint param_2 = 1024u; uint param_2 = 1024u;
Alloc cmd_alloc = slice_mem(param, param_1, param_2); Alloc cmd_alloc = slice_mem(param, param_1, param_2);
CmdRef _1243 = { cmd_alloc.offset }; CmdRef _1357 = { cmd_alloc.offset };
CmdRef cmd_ref = _1243; CmdRef cmd_ref = _1357;
uint cmd_limit = (cmd_ref.offset + 1024u) - 60u; uint cmd_limit = (cmd_ref.offset + 1024u) - 60u;
uint clip_depth = 0u; uint clip_depth = 0u;
uint clip_zero_depth = 0u; uint clip_zero_depth = 0u;
@ -765,17 +832,17 @@ void comp_main()
uint wr_ix = 0u; uint wr_ix = 0u;
uint part_start_ix = 0u; uint part_start_ix = 0u;
uint ready_ix = 0u; uint ready_ix = 0u;
bool mem_ok = _283.Load(4) == 0u; bool mem_ok = _308.Load(4) == 0u;
Alloc param_3; Alloc param_3;
Alloc param_5; Alloc param_5;
uint _1448; uint _1562;
uint element_ix; uint element_ix;
AnnotatedRef ref; AnnotatedRef ref;
Alloc param_14; Alloc param_14;
Alloc param_16; Alloc param_16;
uint tile_count; uint tile_count;
Alloc param_23; Alloc param_23;
uint _1770; uint _1887;
Alloc param_29; Alloc param_29;
Tile tile_1; Tile tile_1;
AnnoColor fill; AnnoColor fill;
@ -783,40 +850,41 @@ void comp_main()
Alloc param_52; Alloc param_52;
CmdLinGrad cmd_lin; CmdLinGrad cmd_lin;
Alloc param_69; Alloc param_69;
Alloc param_95;
while (true) while (true)
{ {
for (uint i = 0u; i < 8u; i++) for (uint i = 0u; i < 8u; i++)
{ {
sh_bitmaps[i][th_ix] = 0u; sh_bitmaps[i][th_ix] = 0u;
} }
bool _1500; bool _1614;
for (;;) for (;;)
{ {
if ((ready_ix == wr_ix) && (partition_ix < n_partitions)) if ((ready_ix == wr_ix) && (partition_ix < n_partitions))
{ {
part_start_ix = ready_ix; part_start_ix = ready_ix;
uint count = 0u; uint count = 0u;
bool _1298 = th_ix < 256u; bool _1412 = th_ix < 256u;
bool _1306; bool _1420;
if (_1298) if (_1412)
{ {
_1306 = (partition_ix + th_ix) < n_partitions; _1420 = (partition_ix + th_ix) < n_partitions;
} }
else else
{ {
_1306 = _1298; _1420 = _1412;
} }
if (_1306) if (_1420)
{ {
uint in_ix = (_1169.Load(20) >> uint(2)) + ((((partition_ix + th_ix) * 256u) + bin_ix) * 2u); uint in_ix = (_1283.Load(20) >> uint(2)) + ((((partition_ix + th_ix) * 256u) + bin_ix) * 2u);
Alloc _1323; Alloc _1437;
_1323.offset = _1169.Load(20); _1437.offset = _1283.Load(20);
param_3.offset = _1323.offset; param_3.offset = _1437.offset;
uint param_4 = in_ix; uint param_4 = in_ix;
count = read_mem(param_3, param_4); count = read_mem(param_3, param_4);
Alloc _1334; Alloc _1448;
_1334.offset = _1169.Load(20); _1448.offset = _1283.Load(20);
param_5.offset = _1334.offset; param_5.offset = _1448.offset;
uint param_6 = in_ix + 1u; uint param_6 = in_ix + 1u;
uint offset = read_mem(param_5, param_6); uint offset = read_mem(param_5, param_6);
uint param_7 = offset; uint param_7 = offset;
@ -862,16 +930,16 @@ void comp_main()
} }
if (part_ix > 0u) if (part_ix > 0u)
{ {
_1448 = sh_part_count[part_ix - 1u]; _1562 = sh_part_count[part_ix - 1u];
} }
else else
{ {
_1448 = part_start_ix; _1562 = part_start_ix;
} }
ix -= _1448; ix -= _1562;
Alloc bin_alloc = sh_part_elements[part_ix]; Alloc bin_alloc = sh_part_elements[part_ix];
BinInstanceRef _1467 = { bin_alloc.offset }; BinInstanceRef _1581 = { bin_alloc.offset };
BinInstanceRef inst_ref = _1467; BinInstanceRef inst_ref = _1581;
BinInstanceRef param_10 = inst_ref; BinInstanceRef param_10 = inst_ref;
uint param_11 = ix; uint param_11 = ix;
Alloc param_12 = bin_alloc; Alloc param_12 = bin_alloc;
@ -881,16 +949,16 @@ void comp_main()
} }
GroupMemoryBarrierWithGroupSync(); GroupMemoryBarrierWithGroupSync();
wr_ix = min((rd_ix + 256u), ready_ix); wr_ix = min((rd_ix + 256u), ready_ix);
bool _1490 = (wr_ix - rd_ix) < 256u; bool _1604 = (wr_ix - rd_ix) < 256u;
if (_1490) if (_1604)
{ {
_1500 = (wr_ix < ready_ix) || (partition_ix < n_partitions); _1614 = (wr_ix < ready_ix) || (partition_ix < n_partitions);
} }
else else
{ {
_1500 = _1490; _1614 = _1604;
} }
if (_1500) if (_1614)
{ {
continue; continue;
} }
@ -903,11 +971,11 @@ void comp_main()
if ((th_ix + rd_ix) < wr_ix) if ((th_ix + rd_ix) < wr_ix)
{ {
element_ix = sh_elements[th_ix]; element_ix = sh_elements[th_ix];
AnnotatedRef _1521 = { _1169.Load(32) + (element_ix * 40u) }; AnnotatedRef _1635 = { _1283.Load(32) + (element_ix * 40u) };
ref = _1521; ref = _1635;
Alloc _1524; Alloc _1638;
_1524.offset = _1169.Load(32); _1638.offset = _1283.Load(32);
param_14.offset = _1524.offset; param_14.offset = _1638.offset;
AnnotatedRef param_15 = ref; AnnotatedRef param_15 = ref;
tag = Annotated_tag(param_14, param_15).tag; tag = Annotated_tag(param_14, param_15).tag;
} }
@ -919,13 +987,13 @@ void comp_main()
case 4u: case 4u:
case 5u: case 5u:
{ {
uint drawmonoid_base = (_1169.Load(44) >> uint(2)) + (2u * element_ix); uint drawmonoid_base = (_1283.Load(44) >> uint(2)) + (2u * element_ix);
uint path_ix = _283.Load(drawmonoid_base * 4 + 8); uint path_ix = _308.Load(drawmonoid_base * 4 + 8);
PathRef _1553 = { _1169.Load(16) + (path_ix * 12u) }; PathRef _1667 = { _1283.Load(16) + (path_ix * 12u) };
Alloc _1556; Alloc _1670;
_1556.offset = _1169.Load(16); _1670.offset = _1283.Load(16);
param_16.offset = _1556.offset; param_16.offset = _1670.offset;
PathRef param_17 = _1553; PathRef param_17 = _1667;
Path path = Path_read(param_16, param_17); Path path = Path_read(param_16, param_17);
uint stride = path.bbox.z - path.bbox.x; uint stride = path.bbox.z - path.bbox.x;
sh_tile_stride[th_ix] = stride; sh_tile_stride[th_ix] = stride;
@ -980,22 +1048,23 @@ void comp_main()
el_ix = probe_1; el_ix = probe_1;
} }
} }
AnnotatedRef _1755 = { _1169.Load(32) + (sh_elements[el_ix] * 40u) }; AnnotatedRef _1869 = { _1283.Load(32) + (sh_elements[el_ix] * 40u) };
AnnotatedRef ref_1 = _1755; AnnotatedRef ref_1 = _1869;
Alloc _1759; Alloc _1874;
_1759.offset = _1169.Load(32); _1874.offset = _1283.Load(32);
param_23.offset = _1759.offset; param_23.offset = _1874.offset;
AnnotatedRef param_24 = ref_1; AnnotatedRef param_24 = ref_1;
uint tag_1 = Annotated_tag(param_23, param_24).tag; AnnotatedTag anno_tag = Annotated_tag(param_23, param_24);
uint tag_1 = anno_tag.tag;
if (el_ix > 0u) if (el_ix > 0u)
{ {
_1770 = sh_tile_count[el_ix - 1u]; _1887 = sh_tile_count[el_ix - 1u];
} }
else else
{ {
_1770 = 0u; _1887 = 0u;
} }
uint seq_ix = ix_1 - _1770; uint seq_ix = ix_1 - _1887;
uint width = sh_tile_width[el_ix]; uint width = sh_tile_width[el_ix];
uint x = sh_tile_x0[el_ix] + (seq_ix % width); uint x = sh_tile_x0[el_ix] + (seq_ix % width);
uint y = sh_tile_y0[el_ix] + (seq_ix / width); uint y = sh_tile_y0[el_ix] + (seq_ix / width);
@ -1004,29 +1073,47 @@ void comp_main()
{ {
uint param_25 = el_ix; uint param_25 = el_ix;
bool param_26 = mem_ok; bool param_26 = mem_ok;
TileRef _1822 = { sh_tile_base[el_ix] + (((sh_tile_stride[el_ix] * y) + x) * 8u) }; TileRef _1939 = { sh_tile_base[el_ix] + (((sh_tile_stride[el_ix] * y) + x) * 8u) };
Alloc param_27 = read_tile_alloc(param_25, param_26); Alloc param_27 = read_tile_alloc(param_25, param_26);
TileRef param_28 = _1822; TileRef param_28 = _1939;
Tile tile = Tile_read(param_27, param_28); Tile tile = Tile_read(param_27, param_28);
bool is_clip = (tag_1 == 4u) || (tag_1 == 5u); bool is_clip = (tag_1 == 4u) || (tag_1 == 5u);
bool _1834 = tile.tile.offset != 0u; bool _1951 = tile.tile.offset != 0u;
bool _1843; bool _1960;
if (!_1834) if (!_1951)
{ {
_1843 = (tile.backdrop == 0) == is_clip; _1960 = (tile.backdrop == 0) == is_clip;
} }
else else
{ {
_1843 = _1834; _1960 = _1951;
} }
include_tile = _1843; bool _1972;
if (!_1960)
{
bool _1971;
if (is_clip)
{
_1971 = (anno_tag.flags & 2u) != 0u;
}
else
{
_1971 = is_clip;
}
_1972 = _1971;
}
else
{
_1972 = _1960;
}
include_tile = _1972;
} }
if (include_tile) if (include_tile)
{ {
uint el_slice = el_ix / 32u; uint el_slice = el_ix / 32u;
uint el_mask = 1u << (el_ix & 31u); uint el_mask = 1u << (el_ix & 31u);
uint _1863; uint _1992;
InterlockedOr(sh_bitmaps[el_slice][(y * 16u) + x], el_mask, _1863); InterlockedOr(sh_bitmaps[el_slice][(y * 16u) + x], el_mask, _1992);
} }
} }
GroupMemoryBarrierWithGroupSync(); GroupMemoryBarrierWithGroupSync();
@ -1050,11 +1137,11 @@ void comp_main()
uint element_ref_ix = (slice_ix * 32u) + uint(int(firstbitlow(bitmap))); uint element_ref_ix = (slice_ix * 32u) + uint(int(firstbitlow(bitmap)));
uint element_ix_1 = sh_elements[element_ref_ix]; uint element_ix_1 = sh_elements[element_ref_ix];
bitmap &= (bitmap - 1u); bitmap &= (bitmap - 1u);
AnnotatedRef _1917 = { _1169.Load(32) + (element_ix_1 * 40u) }; AnnotatedRef _2046 = { _1283.Load(32) + (element_ix_1 * 40u) };
ref = _1917; ref = _2046;
Alloc _1922; Alloc _2050;
_1922.offset = _1169.Load(32); _2050.offset = _1283.Load(32);
param_29.offset = _1922.offset; param_29.offset = _2050.offset;
AnnotatedRef param_30 = ref; AnnotatedRef param_30 = ref;
AnnotatedTag tag_2 = Annotated_tag(param_29, param_30); AnnotatedTag tag_2 = Annotated_tag(param_29, param_30);
if (clip_zero_depth == 0u) if (clip_zero_depth == 0u)
@ -1065,23 +1152,23 @@ void comp_main()
{ {
uint param_31 = element_ref_ix; uint param_31 = element_ref_ix;
bool param_32 = mem_ok; bool param_32 = mem_ok;
TileRef _1958 = { sh_tile_base[element_ref_ix] + (((sh_tile_stride[element_ref_ix] * tile_y) + tile_x) * 8u) }; TileRef _2086 = { sh_tile_base[element_ref_ix] + (((sh_tile_stride[element_ref_ix] * tile_y) + tile_x) * 8u) };
Alloc param_33 = read_tile_alloc(param_31, param_32); Alloc param_33 = read_tile_alloc(param_31, param_32);
TileRef param_34 = _1958; TileRef param_34 = _2086;
tile_1 = Tile_read(param_33, param_34); tile_1 = Tile_read(param_33, param_34);
Alloc _1965; Alloc _2093;
_1965.offset = _1169.Load(32); _2093.offset = _1283.Load(32);
param_35.offset = _1965.offset; param_35.offset = _2093.offset;
AnnotatedRef param_36 = ref; AnnotatedRef param_36 = ref;
fill = Annotated_Color_read(param_35, param_36); fill = Annotated_Color_read(param_35, param_36);
Alloc param_37 = cmd_alloc; Alloc param_37 = cmd_alloc;
CmdRef param_38 = cmd_ref; CmdRef param_38 = cmd_ref;
uint param_39 = cmd_limit; uint param_39 = cmd_limit;
bool _1977 = alloc_cmd(param_37, param_38, param_39); bool _2105 = alloc_cmd(param_37, param_38, param_39);
cmd_alloc = param_37; cmd_alloc = param_37;
cmd_ref = param_38; cmd_ref = param_38;
cmd_limit = param_39; cmd_limit = param_39;
if (!_1977) if (!_2105)
{ {
break; break;
} }
@ -1092,10 +1179,10 @@ void comp_main()
float param_44 = fill.linewidth; float param_44 = fill.linewidth;
write_fill(param_40, param_41, param_42, param_43, param_44); write_fill(param_40, param_41, param_42, param_43, param_44);
cmd_ref = param_41; cmd_ref = param_41;
CmdColor _2001 = { fill.rgba_color }; CmdColor _2129 = { fill.rgba_color };
Alloc param_45 = cmd_alloc; Alloc param_45 = cmd_alloc;
CmdRef param_46 = cmd_ref; CmdRef param_46 = cmd_ref;
CmdColor param_47 = _2001; CmdColor param_47 = _2129;
Cmd_Color_write(param_45, param_46, param_47); Cmd_Color_write(param_45, param_46, param_47);
cmd_ref.offset += 8u; cmd_ref.offset += 8u;
break; break;
@ -1104,23 +1191,23 @@ void comp_main()
{ {
uint param_48 = element_ref_ix; uint param_48 = element_ref_ix;
bool param_49 = mem_ok; bool param_49 = mem_ok;
TileRef _2030 = { sh_tile_base[element_ref_ix] + (((sh_tile_stride[element_ref_ix] * tile_y) + tile_x) * 8u) }; TileRef _2158 = { sh_tile_base[element_ref_ix] + (((sh_tile_stride[element_ref_ix] * tile_y) + tile_x) * 8u) };
Alloc param_50 = read_tile_alloc(param_48, param_49); Alloc param_50 = read_tile_alloc(param_48, param_49);
TileRef param_51 = _2030; TileRef param_51 = _2158;
tile_1 = Tile_read(param_50, param_51); tile_1 = Tile_read(param_50, param_51);
Alloc _2037; Alloc _2165;
_2037.offset = _1169.Load(32); _2165.offset = _1283.Load(32);
param_52.offset = _2037.offset; param_52.offset = _2165.offset;
AnnotatedRef param_53 = ref; AnnotatedRef param_53 = ref;
AnnoLinGradient lin = Annotated_LinGradient_read(param_52, param_53); AnnoLinGradient lin = Annotated_LinGradient_read(param_52, param_53);
Alloc param_54 = cmd_alloc; Alloc param_54 = cmd_alloc;
CmdRef param_55 = cmd_ref; CmdRef param_55 = cmd_ref;
uint param_56 = cmd_limit; uint param_56 = cmd_limit;
bool _2049 = alloc_cmd(param_54, param_55, param_56); bool _2177 = alloc_cmd(param_54, param_55, param_56);
cmd_alloc = param_54; cmd_alloc = param_54;
cmd_ref = param_55; cmd_ref = param_55;
cmd_limit = param_56; cmd_limit = param_56;
if (!_2049) if (!_2177)
{ {
break; break;
} }
@ -1146,23 +1233,23 @@ void comp_main()
{ {
uint param_65 = element_ref_ix; uint param_65 = element_ref_ix;
bool param_66 = mem_ok; bool param_66 = mem_ok;
TileRef _2114 = { sh_tile_base[element_ref_ix] + (((sh_tile_stride[element_ref_ix] * tile_y) + tile_x) * 8u) }; TileRef _2242 = { sh_tile_base[element_ref_ix] + (((sh_tile_stride[element_ref_ix] * tile_y) + tile_x) * 8u) };
Alloc param_67 = read_tile_alloc(param_65, param_66); Alloc param_67 = read_tile_alloc(param_65, param_66);
TileRef param_68 = _2114; TileRef param_68 = _2242;
tile_1 = Tile_read(param_67, param_68); tile_1 = Tile_read(param_67, param_68);
Alloc _2121; Alloc _2249;
_2121.offset = _1169.Load(32); _2249.offset = _1283.Load(32);
param_69.offset = _2121.offset; param_69.offset = _2249.offset;
AnnotatedRef param_70 = ref; AnnotatedRef param_70 = ref;
AnnoImage fill_img = Annotated_Image_read(param_69, param_70); AnnoImage fill_img = Annotated_Image_read(param_69, param_70);
Alloc param_71 = cmd_alloc; Alloc param_71 = cmd_alloc;
CmdRef param_72 = cmd_ref; CmdRef param_72 = cmd_ref;
uint param_73 = cmd_limit; uint param_73 = cmd_limit;
bool _2133 = alloc_cmd(param_71, param_72, param_73); bool _2261 = alloc_cmd(param_71, param_72, param_73);
cmd_alloc = param_71; cmd_alloc = param_71;
cmd_ref = param_72; cmd_ref = param_72;
cmd_limit = param_73; cmd_limit = param_73;
if (!_2133) if (!_2261)
{ {
break; break;
} }
@ -1173,10 +1260,10 @@ void comp_main()
float param_78 = fill_img.linewidth; float param_78 = fill_img.linewidth;
write_fill(param_74, param_75, param_76, param_77, param_78); write_fill(param_74, param_75, param_76, param_77, param_78);
cmd_ref = param_75; cmd_ref = param_75;
CmdImage _2159 = { fill_img.index, fill_img.offset }; CmdImage _2287 = { fill_img.index, fill_img.offset };
Alloc param_79 = cmd_alloc; Alloc param_79 = cmd_alloc;
CmdRef param_80 = cmd_ref; CmdRef param_80 = cmd_ref;
CmdImage param_81 = _2159; CmdImage param_81 = _2287;
Cmd_Image_write(param_79, param_80, param_81); Cmd_Image_write(param_79, param_80, param_81);
cmd_ref.offset += 12u; cmd_ref.offset += 12u;
break; break;
@ -1185,21 +1272,21 @@ void comp_main()
{ {
uint param_82 = element_ref_ix; uint param_82 = element_ref_ix;
bool param_83 = mem_ok; bool param_83 = mem_ok;
TileRef _2188 = { sh_tile_base[element_ref_ix] + (((sh_tile_stride[element_ref_ix] * tile_y) + tile_x) * 8u) }; TileRef _2316 = { sh_tile_base[element_ref_ix] + (((sh_tile_stride[element_ref_ix] * tile_y) + tile_x) * 8u) };
Alloc param_84 = read_tile_alloc(param_82, param_83); Alloc param_84 = read_tile_alloc(param_82, param_83);
TileRef param_85 = _2188; TileRef param_85 = _2316;
tile_1 = Tile_read(param_84, param_85); tile_1 = Tile_read(param_84, param_85);
bool _2194 = tile_1.tile.offset == 0u; bool _2322 = tile_1.tile.offset == 0u;
bool _2200; bool _2328;
if (_2194) if (_2322)
{ {
_2200 = tile_1.backdrop == 0; _2328 = tile_1.backdrop == 0;
} }
else else
{ {
_2200 = _2194; _2328 = _2322;
} }
if (_2200) if (_2328)
{ {
clip_zero_depth = clip_depth + 1u; clip_zero_depth = clip_depth + 1u;
} }
@ -1208,11 +1295,11 @@ void comp_main()
Alloc param_86 = cmd_alloc; Alloc param_86 = cmd_alloc;
CmdRef param_87 = cmd_ref; CmdRef param_87 = cmd_ref;
uint param_88 = cmd_limit; uint param_88 = cmd_limit;
bool _2212 = alloc_cmd(param_86, param_87, param_88); bool _2340 = alloc_cmd(param_86, param_87, param_88);
cmd_alloc = param_86; cmd_alloc = param_86;
cmd_ref = param_87; cmd_ref = param_87;
cmd_limit = param_88; cmd_limit = param_88;
if (!_2212) if (!_2340)
{ {
break; break;
} }
@ -1228,33 +1315,40 @@ void comp_main()
{ {
uint param_91 = element_ref_ix; uint param_91 = element_ref_ix;
bool param_92 = mem_ok; bool param_92 = mem_ok;
TileRef _2249 = { sh_tile_base[element_ref_ix] + (((sh_tile_stride[element_ref_ix] * tile_y) + tile_x) * 8u) }; TileRef _2377 = { sh_tile_base[element_ref_ix] + (((sh_tile_stride[element_ref_ix] * tile_y) + tile_x) * 8u) };
Alloc param_93 = read_tile_alloc(param_91, param_92); Alloc param_93 = read_tile_alloc(param_91, param_92);
TileRef param_94 = _2249; TileRef param_94 = _2377;
tile_1 = Tile_read(param_93, param_94); tile_1 = Tile_read(param_93, param_94);
Alloc _2384;
_2384.offset = _1283.Load(32);
param_95.offset = _2384.offset;
AnnotatedRef param_96 = ref;
AnnoEndClip end_clip = Annotated_EndClip_read(param_95, param_96);
clip_depth--; clip_depth--;
Alloc param_95 = cmd_alloc; Alloc param_97 = cmd_alloc;
CmdRef param_96 = cmd_ref; CmdRef param_98 = cmd_ref;
uint param_97 = cmd_limit; uint param_99 = cmd_limit;
bool _2261 = alloc_cmd(param_95, param_96, param_97); bool _2398 = alloc_cmd(param_97, param_98, param_99);
cmd_alloc = param_95; cmd_alloc = param_97;
cmd_ref = param_96; cmd_ref = param_98;
cmd_limit = param_97; cmd_limit = param_99;
if (!_2261) if (!_2398)
{ {
break; break;
} }
Alloc param_98 = cmd_alloc; Alloc param_100 = cmd_alloc;
CmdRef param_99 = cmd_ref; CmdRef param_101 = cmd_ref;
uint param_100 = 0u; uint param_102 = 0u;
Tile param_101 = tile_1; Tile param_103 = tile_1;
float param_102 = 0.0f; float param_104 = 0.0f;
write_fill(param_98, param_99, param_100, param_101, param_102); write_fill(param_100, param_101, param_102, param_103, param_104);
cmd_ref = param_99; cmd_ref = param_101;
Alloc param_103 = cmd_alloc; CmdEndClip _2419 = { end_clip.blend };
CmdRef param_104 = cmd_ref; Alloc param_105 = cmd_alloc;
Cmd_EndClip_write(param_103, param_104); CmdRef param_106 = cmd_ref;
cmd_ref.offset += 4u; CmdEndClip param_107 = _2419;
Cmd_EndClip_write(param_105, param_106, param_107);
cmd_ref.offset += 8u;
break; break;
} }
} }
@ -1287,21 +1381,21 @@ void comp_main()
break; break;
} }
} }
bool _2326 = (bin_tile_x + tile_x) < _1169.Load(8); bool _2467 = (bin_tile_x + tile_x) < _1283.Load(8);
bool _2335; bool _2476;
if (_2326) if (_2467)
{ {
_2335 = (bin_tile_y + tile_y) < _1169.Load(12); _2476 = (bin_tile_y + tile_y) < _1283.Load(12);
} }
else else
{ {
_2335 = _2326; _2476 = _2467;
} }
if (_2335) if (_2476)
{ {
Alloc param_105 = cmd_alloc; Alloc param_108 = cmd_alloc;
CmdRef param_106 = cmd_ref; CmdRef param_109 = cmd_ref;
Cmd_End_write(param_105, param_106); Cmd_End_write(param_108, param_109);
} }
} }

14
piet-gpu/shader/gen/coarse.msl generated
View file

@ -7,6 +7,13 @@
using namespace metal; using namespace metal;
// Implementation of the GLSL findLSB() function
template<typename T>
inline T spvFindLSB(T x)
{
return select(ctz(x), T(-1), x == T(0));
}
struct Alloc struct Alloc
{ {
uint offset; uint offset;
@ -244,13 +251,6 @@ struct ConfigBuf
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(256u, 1u, 1u); constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(256u, 1u, 1u);
// Implementation of the GLSL findLSB() function
template<typename T>
inline T spvFindLSB(T x)
{
return select(ctz(x), T(-1), x == T(0));
}
static inline __attribute__((always_inline)) static inline __attribute__((always_inline))
Alloc slice_mem(thread const Alloc& a, thread const uint& offset, thread const uint& size) Alloc slice_mem(thread const Alloc& a, thread const uint& offset, thread const uint& size)
{ {

BIN
piet-gpu/shader/gen/draw_leaf.dxil generated
View file

Binary file not shown.

232
piet-gpu/shader/gen/draw_leaf.hlsl generated
View file

@ -41,6 +41,17 @@ struct FillImage
int2 offset; int2 offset;
}; };
struct ClipRef
{
uint offset;
};
struct Clip
{
float4 bbox;
uint blend;
};
struct ElementTag struct ElementTag
{ {
uint tag; uint tag;
@ -102,6 +113,7 @@ struct AnnoBeginClip
{ {
float4 bbox; float4 bbox;
float linewidth; float linewidth;
uint blend;
}; };
struct AnnoEndClipRef struct AnnoEndClipRef
@ -112,6 +124,7 @@ struct AnnoEndClipRef
struct AnnoEndClip struct AnnoEndClip
{ {
float4 bbox; float4 bbox;
uint blend;
}; };
struct AnnotatedRef struct AnnotatedRef
@ -148,14 +161,14 @@ struct Config
static const uint3 gl_WorkGroupSize = uint3(256u, 1u, 1u); static const uint3 gl_WorkGroupSize = uint3(256u, 1u, 1u);
static const DrawMonoid _348 = { 0u, 0u }; static const DrawMonoid _413 = { 0u, 0u };
static const DrawMonoid _372 = { 1u, 0u }; static const DrawMonoid _437 = { 1u, 0u };
static const DrawMonoid _374 = { 1u, 1u }; static const DrawMonoid _439 = { 1u, 1u };
RWByteAddressBuffer _187 : register(u0, space0); RWByteAddressBuffer _199 : register(u0, space0);
ByteAddressBuffer _211 : register(t2, space0); ByteAddressBuffer _223 : register(t2, space0);
ByteAddressBuffer _934 : register(t3, space0); ByteAddressBuffer _1020 : register(t3, space0);
ByteAddressBuffer _968 : register(t1, space0); ByteAddressBuffer _1054 : register(t1, space0);
static uint3 gl_WorkGroupID; static uint3 gl_WorkGroupID;
static uint3 gl_LocalInvocationID; static uint3 gl_LocalInvocationID;
@ -171,9 +184,9 @@ groupshared DrawMonoid sh_scratch[256];
ElementTag Element_tag(ElementRef ref) ElementTag Element_tag(ElementRef ref)
{ {
uint tag_and_flags = _211.Load((ref.offset >> uint(2)) * 4 + 0); uint tag_and_flags = _223.Load((ref.offset >> uint(2)) * 4 + 0);
ElementTag _321 = { tag_and_flags & 65535u, tag_and_flags >> uint(16) }; ElementTag _378 = { tag_and_flags & 65535u, tag_and_flags >> uint(16) };
return _321; return _378;
} }
DrawMonoid map_tag(uint tag_word) DrawMonoid map_tag(uint tag_word)
@ -184,24 +197,24 @@ DrawMonoid map_tag(uint tag_word)
case 5u: case 5u:
case 6u: case 6u:
{ {
return _372; return _437;
} }
case 9u: case 9u:
case 10u: case 10u:
{ {
return _374; return _439;
} }
default: default:
{ {
return _348; return _413;
} }
} }
} }
ElementRef Element_index(ElementRef ref, uint index) ElementRef Element_index(ElementRef ref, uint index)
{ {
ElementRef _200 = { ref.offset + (index * 36u) }; ElementRef _212 = { ref.offset + (index * 36u) };
return _200; return _212;
} }
DrawMonoid combine_tag_monoid(DrawMonoid a, DrawMonoid b) DrawMonoid combine_tag_monoid(DrawMonoid a, DrawMonoid b)
@ -214,13 +227,13 @@ DrawMonoid combine_tag_monoid(DrawMonoid a, DrawMonoid b)
DrawMonoid tag_monoid_identity() DrawMonoid tag_monoid_identity()
{ {
return _348; return _413;
} }
FillColor FillColor_read(FillColorRef ref) FillColor FillColor_read(FillColorRef ref)
{ {
uint ix = ref.offset >> uint(2); uint ix = ref.offset >> uint(2);
uint raw0 = _211.Load((ix + 0u) * 4 + 0); uint raw0 = _223.Load((ix + 0u) * 4 + 0);
FillColor s; FillColor s;
s.rgba_color = raw0; s.rgba_color = raw0;
return s; return s;
@ -228,8 +241,8 @@ FillColor FillColor_read(FillColorRef ref)
FillColor Element_FillColor_read(ElementRef ref) FillColor Element_FillColor_read(ElementRef ref)
{ {
FillColorRef _327 = { ref.offset + 4u }; FillColorRef _384 = { ref.offset + 4u };
FillColorRef param = _327; FillColorRef param = _384;
return FillColor_read(param); return FillColor_read(param);
} }
@ -246,7 +259,7 @@ void write_mem(Alloc alloc, uint offset, uint val)
{ {
return; return;
} }
_187.Store(offset * 4 + 8, val); _199.Store(offset * 4 + 8, val);
} }
void AnnoColor_write(Alloc a, AnnoColorRef ref, AnnoColor s) void AnnoColor_write(Alloc a, AnnoColorRef ref, AnnoColor s)
@ -284,9 +297,9 @@ void Annotated_Color_write(Alloc a, AnnotatedRef ref, uint flags, AnnoColor s)
uint param_1 = ref.offset >> uint(2); uint param_1 = ref.offset >> uint(2);
uint param_2 = (flags << uint(16)) | 1u; uint param_2 = (flags << uint(16)) | 1u;
write_mem(param, param_1, param_2); write_mem(param, param_1, param_2);
AnnoColorRef _735 = { ref.offset + 4u }; AnnoColorRef _818 = { ref.offset + 4u };
Alloc param_3 = a; Alloc param_3 = a;
AnnoColorRef param_4 = _735; AnnoColorRef param_4 = _818;
AnnoColor param_5 = s; AnnoColor param_5 = s;
AnnoColor_write(param_3, param_4, param_5); AnnoColor_write(param_3, param_4, param_5);
} }
@ -294,11 +307,11 @@ void Annotated_Color_write(Alloc a, AnnotatedRef ref, uint flags, AnnoColor s)
FillLinGradient FillLinGradient_read(FillLinGradientRef ref) FillLinGradient FillLinGradient_read(FillLinGradientRef ref)
{ {
uint ix = ref.offset >> uint(2); uint ix = ref.offset >> uint(2);
uint raw0 = _211.Load((ix + 0u) * 4 + 0); uint raw0 = _223.Load((ix + 0u) * 4 + 0);
uint raw1 = _211.Load((ix + 1u) * 4 + 0); uint raw1 = _223.Load((ix + 1u) * 4 + 0);
uint raw2 = _211.Load((ix + 2u) * 4 + 0); uint raw2 = _223.Load((ix + 2u) * 4 + 0);
uint raw3 = _211.Load((ix + 3u) * 4 + 0); uint raw3 = _223.Load((ix + 3u) * 4 + 0);
uint raw4 = _211.Load((ix + 4u) * 4 + 0); uint raw4 = _223.Load((ix + 4u) * 4 + 0);
FillLinGradient s; FillLinGradient s;
s.index = raw0; s.index = raw0;
s.p0 = float2(asfloat(raw1), asfloat(raw2)); s.p0 = float2(asfloat(raw1), asfloat(raw2));
@ -308,8 +321,8 @@ FillLinGradient FillLinGradient_read(FillLinGradientRef ref)
FillLinGradient Element_FillLinGradient_read(ElementRef ref) FillLinGradient Element_FillLinGradient_read(ElementRef ref)
{ {
FillLinGradientRef _335 = { ref.offset + 4u }; FillLinGradientRef _392 = { ref.offset + 4u };
FillLinGradientRef param = _335; FillLinGradientRef param = _392;
return FillLinGradient_read(param); return FillLinGradient_read(param);
} }
@ -360,9 +373,9 @@ void Annotated_LinGradient_write(Alloc a, AnnotatedRef ref, uint flags, AnnoLinG
uint param_1 = ref.offset >> uint(2); uint param_1 = ref.offset >> uint(2);
uint param_2 = (flags << uint(16)) | 2u; uint param_2 = (flags << uint(16)) | 2u;
write_mem(param, param_1, param_2); write_mem(param, param_1, param_2);
AnnoLinGradientRef _756 = { ref.offset + 4u }; AnnoLinGradientRef _839 = { ref.offset + 4u };
Alloc param_3 = a; Alloc param_3 = a;
AnnoLinGradientRef param_4 = _756; AnnoLinGradientRef param_4 = _839;
AnnoLinGradient param_5 = s; AnnoLinGradient param_5 = s;
AnnoLinGradient_write(param_3, param_4, param_5); AnnoLinGradient_write(param_3, param_4, param_5);
} }
@ -370,8 +383,8 @@ void Annotated_LinGradient_write(Alloc a, AnnotatedRef ref, uint flags, AnnoLinG
FillImage FillImage_read(FillImageRef ref) FillImage FillImage_read(FillImageRef ref)
{ {
uint ix = ref.offset >> uint(2); uint ix = ref.offset >> uint(2);
uint raw0 = _211.Load((ix + 0u) * 4 + 0); uint raw0 = _223.Load((ix + 0u) * 4 + 0);
uint raw1 = _211.Load((ix + 1u) * 4 + 0); uint raw1 = _223.Load((ix + 1u) * 4 + 0);
FillImage s; FillImage s;
s.index = raw0; s.index = raw0;
s.offset = int2(int(raw1 << uint(16)) >> 16, int(raw1) >> 16); s.offset = int2(int(raw1 << uint(16)) >> 16, int(raw1) >> 16);
@ -380,8 +393,8 @@ FillImage FillImage_read(FillImageRef ref)
FillImage Element_FillImage_read(ElementRef ref) FillImage Element_FillImage_read(ElementRef ref)
{ {
FillImageRef _343 = { ref.offset + 4u }; FillImageRef _400 = { ref.offset + 4u };
FillImageRef param = _343; FillImageRef param = _400;
return FillImage_read(param); return FillImage_read(param);
} }
@ -424,13 +437,33 @@ void Annotated_Image_write(Alloc a, AnnotatedRef ref, uint flags, AnnoImage s)
uint param_1 = ref.offset >> uint(2); uint param_1 = ref.offset >> uint(2);
uint param_2 = (flags << uint(16)) | 3u; uint param_2 = (flags << uint(16)) | 3u;
write_mem(param, param_1, param_2); write_mem(param, param_1, param_2);
AnnoImageRef _777 = { ref.offset + 4u }; AnnoImageRef _860 = { ref.offset + 4u };
Alloc param_3 = a; Alloc param_3 = a;
AnnoImageRef param_4 = _777; AnnoImageRef param_4 = _860;
AnnoImage param_5 = s; AnnoImage param_5 = s;
AnnoImage_write(param_3, param_4, param_5); AnnoImage_write(param_3, param_4, param_5);
} }
Clip Clip_read(ClipRef ref)
{
uint ix = ref.offset >> uint(2);
uint raw0 = _223.Load((ix + 0u) * 4 + 0);
uint raw1 = _223.Load((ix + 1u) * 4 + 0);
uint raw2 = _223.Load((ix + 2u) * 4 + 0);
uint raw3 = _223.Load((ix + 3u) * 4 + 0);
Clip s;
s.bbox = float4(asfloat(raw0), asfloat(raw1), asfloat(raw2), asfloat(raw3));
s.blend = _223.Load((ix + 4u) * 4 + 0);
return s;
}
Clip Element_BeginClip_read(ElementRef ref)
{
ClipRef _408 = { ref.offset + 4u };
ClipRef param = _408;
return Clip_read(param);
}
void AnnoBeginClip_write(Alloc a, AnnoBeginClipRef ref, AnnoBeginClip s) void AnnoBeginClip_write(Alloc a, AnnoBeginClipRef ref, AnnoBeginClip s)
{ {
uint ix = ref.offset >> uint(2); uint ix = ref.offset >> uint(2);
@ -454,6 +487,10 @@ void AnnoBeginClip_write(Alloc a, AnnoBeginClipRef ref, AnnoBeginClip s)
uint param_13 = ix + 4u; uint param_13 = ix + 4u;
uint param_14 = asuint(s.linewidth); uint param_14 = asuint(s.linewidth);
write_mem(param_12, param_13, param_14); write_mem(param_12, param_13, param_14);
Alloc param_15 = a;
uint param_16 = ix + 5u;
uint param_17 = s.blend;
write_mem(param_15, param_16, param_17);
} }
void Annotated_BeginClip_write(Alloc a, AnnotatedRef ref, uint flags, AnnoBeginClip s) void Annotated_BeginClip_write(Alloc a, AnnotatedRef ref, uint flags, AnnoBeginClip s)
@ -462,9 +499,9 @@ void Annotated_BeginClip_write(Alloc a, AnnotatedRef ref, uint flags, AnnoBeginC
uint param_1 = ref.offset >> uint(2); uint param_1 = ref.offset >> uint(2);
uint param_2 = (flags << uint(16)) | 4u; uint param_2 = (flags << uint(16)) | 4u;
write_mem(param, param_1, param_2); write_mem(param, param_1, param_2);
AnnoBeginClipRef _798 = { ref.offset + 4u }; AnnoBeginClipRef _881 = { ref.offset + 4u };
Alloc param_3 = a; Alloc param_3 = a;
AnnoBeginClipRef param_4 = _798; AnnoBeginClipRef param_4 = _881;
AnnoBeginClip param_5 = s; AnnoBeginClip param_5 = s;
AnnoBeginClip_write(param_3, param_4, param_5); AnnoBeginClip_write(param_3, param_4, param_5);
} }
@ -488,17 +525,21 @@ void AnnoEndClip_write(Alloc a, AnnoEndClipRef ref, AnnoEndClip s)
uint param_10 = ix + 3u; uint param_10 = ix + 3u;
uint param_11 = asuint(s.bbox.w); uint param_11 = asuint(s.bbox.w);
write_mem(param_9, param_10, param_11); write_mem(param_9, param_10, param_11);
Alloc param_12 = a;
uint param_13 = ix + 4u;
uint param_14 = s.blend;
write_mem(param_12, param_13, param_14);
} }
void Annotated_EndClip_write(Alloc a, AnnotatedRef ref, AnnoEndClip s) void Annotated_EndClip_write(Alloc a, AnnotatedRef ref, uint flags, AnnoEndClip s)
{ {
Alloc param = a; Alloc param = a;
uint param_1 = ref.offset >> uint(2); uint param_1 = ref.offset >> uint(2);
uint param_2 = 5u; uint param_2 = (flags << uint(16)) | 5u;
write_mem(param, param_1, param_2); write_mem(param, param_1, param_2);
AnnoEndClipRef _816 = { ref.offset + 4u }; AnnoEndClipRef _902 = { ref.offset + 4u };
Alloc param_3 = a; Alloc param_3 = a;
AnnoEndClipRef param_4 = _816; AnnoEndClipRef param_4 = _902;
AnnoEndClip param_5 = s; AnnoEndClip param_5 = s;
AnnoEndClip_write(param_3, param_4, param_5); AnnoEndClip_write(param_3, param_4, param_5);
} }
@ -506,8 +547,8 @@ void Annotated_EndClip_write(Alloc a, AnnotatedRef ref, AnnoEndClip s)
void comp_main() void comp_main()
{ {
uint ix = gl_GlobalInvocationID.x * 8u; uint ix = gl_GlobalInvocationID.x * 8u;
ElementRef _834 = { ix * 36u }; ElementRef _920 = { ix * 36u };
ElementRef ref = _834; ElementRef ref = _920;
ElementRef param = ref; ElementRef param = ref;
uint tag_word = Element_tag(param).tag; uint tag_word = Element_tag(param).tag;
uint param_1 = tag_word; uint param_1 = tag_word;
@ -544,11 +585,11 @@ void comp_main()
DrawMonoid row = tag_monoid_identity(); DrawMonoid row = tag_monoid_identity();
if (gl_WorkGroupID.x > 0u) if (gl_WorkGroupID.x > 0u)
{ {
DrawMonoid _940; DrawMonoid _1026;
_940.path_ix = _934.Load((gl_WorkGroupID.x - 1u) * 8 + 0); _1026.path_ix = _1020.Load((gl_WorkGroupID.x - 1u) * 8 + 0);
_940.clip_ix = _934.Load((gl_WorkGroupID.x - 1u) * 8 + 4); _1026.clip_ix = _1020.Load((gl_WorkGroupID.x - 1u) * 8 + 4);
row.path_ix = _940.path_ix; row.path_ix = _1026.path_ix;
row.clip_ix = _940.clip_ix; row.clip_ix = _1026.clip_ix;
} }
if (gl_LocalInvocationID.x > 0u) if (gl_LocalInvocationID.x > 0u)
{ {
@ -557,10 +598,10 @@ void comp_main()
row = combine_tag_monoid(param_10, param_11); row = combine_tag_monoid(param_10, param_11);
} }
uint out_ix = gl_GlobalInvocationID.x * 8u; uint out_ix = gl_GlobalInvocationID.x * 8u;
uint out_base = (_968.Load(44) >> uint(2)) + (out_ix * 2u); uint out_base = (_1054.Load(44) >> uint(2)) + (out_ix * 2u);
uint clip_out_base = _968.Load(48) >> uint(2); uint clip_out_base = _1054.Load(48) >> uint(2);
AnnotatedRef _989 = { _968.Load(32) + (out_ix * 40u) }; AnnotatedRef _1075 = { _1054.Load(32) + (out_ix * 40u) };
AnnotatedRef out_ref = _989; AnnotatedRef out_ref = _1075;
float4 mat; float4 mat;
float2 translate; float2 translate;
AnnoColor anno_fill; AnnoColor anno_fill;
@ -570,9 +611,9 @@ void comp_main()
AnnoImage anno_img; AnnoImage anno_img;
Alloc param_28; Alloc param_28;
AnnoBeginClip anno_begin_clip; AnnoBeginClip anno_begin_clip;
Alloc param_32; Alloc param_33;
AnnoEndClip anno_end_clip; AnnoEndClip anno_end_clip;
Alloc param_36; Alloc param_38;
for (uint i_2 = 0u; i_2 < 8u; i_2++) for (uint i_2 = 0u; i_2 < 8u; i_2++)
{ {
DrawMonoid m = row; DrawMonoid m = row;
@ -582,8 +623,8 @@ void comp_main()
DrawMonoid param_13 = local[i_2 - 1u]; DrawMonoid param_13 = local[i_2 - 1u];
m = combine_tag_monoid(param_12, param_13); m = combine_tag_monoid(param_12, param_13);
} }
_187.Store((out_base + (i_2 * 2u)) * 4 + 8, m.path_ix); _199.Store((out_base + (i_2 * 2u)) * 4 + 8, m.path_ix);
_187.Store(((out_base + (i_2 * 2u)) + 1u) * 4 + 8, m.clip_ix); _199.Store(((out_base + (i_2 * 2u)) + 1u) * 4 + 8, m.clip_ix);
ElementRef param_14 = ref; ElementRef param_14 = ref;
uint param_15 = i_2; uint param_15 = i_2;
ElementRef this_ref = Element_index(param_14, param_15); ElementRef this_ref = Element_index(param_14, param_15);
@ -591,22 +632,22 @@ void comp_main()
tag_word = Element_tag(param_16).tag; tag_word = Element_tag(param_16).tag;
if ((((tag_word == 4u) || (tag_word == 5u)) || (tag_word == 6u)) || (tag_word == 9u)) if ((((tag_word == 4u) || (tag_word == 5u)) || (tag_word == 6u)) || (tag_word == 9u))
{ {
uint bbox_offset = (_968.Load(40) >> uint(2)) + (6u * m.path_ix); uint bbox_offset = (_1054.Load(40) >> uint(2)) + (6u * m.path_ix);
float bbox_l = float(_187.Load(bbox_offset * 4 + 8)) - 32768.0f; float bbox_l = float(_199.Load(bbox_offset * 4 + 8)) - 32768.0f;
float bbox_t = float(_187.Load((bbox_offset + 1u) * 4 + 8)) - 32768.0f; float bbox_t = float(_199.Load((bbox_offset + 1u) * 4 + 8)) - 32768.0f;
float bbox_r = float(_187.Load((bbox_offset + 2u) * 4 + 8)) - 32768.0f; float bbox_r = float(_199.Load((bbox_offset + 2u) * 4 + 8)) - 32768.0f;
float bbox_b = float(_187.Load((bbox_offset + 3u) * 4 + 8)) - 32768.0f; float bbox_b = float(_199.Load((bbox_offset + 3u) * 4 + 8)) - 32768.0f;
float4 bbox = float4(bbox_l, bbox_t, bbox_r, bbox_b); float4 bbox = float4(bbox_l, bbox_t, bbox_r, bbox_b);
float linewidth = asfloat(_187.Load((bbox_offset + 4u) * 4 + 8)); float linewidth = asfloat(_199.Load((bbox_offset + 4u) * 4 + 8));
uint fill_mode = uint(linewidth >= 0.0f); uint fill_mode = uint(linewidth >= 0.0f);
if ((linewidth >= 0.0f) || (tag_word == 5u)) if ((linewidth >= 0.0f) || (tag_word == 5u))
{ {
uint trans_ix = _187.Load((bbox_offset + 5u) * 4 + 8); uint trans_ix = _199.Load((bbox_offset + 5u) * 4 + 8);
uint t = (_968.Load(36) >> uint(2)) + (6u * trans_ix); uint t = (_1054.Load(36) >> uint(2)) + (6u * trans_ix);
mat = asfloat(uint4(_187.Load(t * 4 + 8), _187.Load((t + 1u) * 4 + 8), _187.Load((t + 2u) * 4 + 8), _187.Load((t + 3u) * 4 + 8))); mat = asfloat(uint4(_199.Load(t * 4 + 8), _199.Load((t + 1u) * 4 + 8), _199.Load((t + 2u) * 4 + 8), _199.Load((t + 3u) * 4 + 8)));
if (tag_word == 5u) if (tag_word == 5u)
{ {
translate = asfloat(uint2(_187.Load((t + 4u) * 4 + 8), _187.Load((t + 5u) * 4 + 8))); translate = asfloat(uint2(_199.Load((t + 4u) * 4 + 8), _199.Load((t + 5u) * 4 + 8)));
} }
} }
if (linewidth >= 0.0f) if (linewidth >= 0.0f)
@ -623,9 +664,9 @@ void comp_main()
anno_fill.bbox = bbox; anno_fill.bbox = bbox;
anno_fill.linewidth = linewidth; anno_fill.linewidth = linewidth;
anno_fill.rgba_color = fill.rgba_color; anno_fill.rgba_color = fill.rgba_color;
Alloc _1203; Alloc _1288;
_1203.offset = _968.Load(32); _1288.offset = _1054.Load(32);
param_18.offset = _1203.offset; param_18.offset = _1288.offset;
AnnotatedRef param_19 = out_ref; AnnotatedRef param_19 = out_ref;
uint param_20 = fill_mode; uint param_20 = fill_mode;
AnnoColor param_21 = anno_fill; AnnoColor param_21 = anno_fill;
@ -648,9 +689,9 @@ void comp_main()
anno_lin.line_x = line_x; anno_lin.line_x = line_x;
anno_lin.line_y = line_y; anno_lin.line_y = line_y;
anno_lin.line_c = -((p0.x * line_x) + (p0.y * line_y)); anno_lin.line_c = -((p0.x * line_x) + (p0.y * line_y));
Alloc _1299; Alloc _1384;
_1299.offset = _968.Load(32); _1384.offset = _1054.Load(32);
param_23.offset = _1299.offset; param_23.offset = _1384.offset;
AnnotatedRef param_24 = out_ref; AnnotatedRef param_24 = out_ref;
uint param_25 = fill_mode; uint param_25 = fill_mode;
AnnoLinGradient param_26 = anno_lin; AnnoLinGradient param_26 = anno_lin;
@ -665,9 +706,9 @@ void comp_main()
anno_img.linewidth = linewidth; anno_img.linewidth = linewidth;
anno_img.index = fill_img.index; anno_img.index = fill_img.index;
anno_img.offset = fill_img.offset; anno_img.offset = fill_img.offset;
Alloc _1327; Alloc _1412;
_1327.offset = _968.Load(32); _1412.offset = _1054.Load(32);
param_28.offset = _1327.offset; param_28.offset = _1412.offset;
AnnotatedRef param_29 = out_ref; AnnotatedRef param_29 = out_ref;
uint param_30 = fill_mode; uint param_30 = fill_mode;
AnnoImage param_31 = anno_img; AnnoImage param_31 = anno_img;
@ -676,15 +717,19 @@ void comp_main()
} }
case 9u: case 9u:
{ {
ElementRef param_32 = this_ref;
Clip begin_clip = Element_BeginClip_read(param_32);
anno_begin_clip.bbox = bbox; anno_begin_clip.bbox = bbox;
anno_begin_clip.linewidth = 0.0f; anno_begin_clip.linewidth = 0.0f;
Alloc _1344; anno_begin_clip.blend = begin_clip.blend;
_1344.offset = _968.Load(32); uint flags = uint(begin_clip.blend != 3u) << uint(1);
param_32.offset = _1344.offset; Alloc _1442;
AnnotatedRef param_33 = out_ref; _1442.offset = _1054.Load(32);
uint param_34 = 0u; param_33.offset = _1442.offset;
AnnoBeginClip param_35 = anno_begin_clip; AnnotatedRef param_34 = out_ref;
Annotated_BeginClip_write(param_32, param_33, param_34, param_35); uint param_35 = flags;
AnnoBeginClip param_36 = anno_begin_clip;
Annotated_BeginClip_write(param_33, param_34, param_35, param_36);
break; break;
} }
} }
@ -693,13 +738,18 @@ void comp_main()
{ {
if (tag_word == 10u) if (tag_word == 10u)
{ {
ElementRef param_37 = this_ref;
Clip end_clip = Element_BeginClip_read(param_37);
anno_end_clip.bbox = float4(-1000000000.0f, -1000000000.0f, 1000000000.0f, 1000000000.0f); anno_end_clip.bbox = float4(-1000000000.0f, -1000000000.0f, 1000000000.0f, 1000000000.0f);
Alloc _1368; anno_end_clip.blend = end_clip.blend;
_1368.offset = _968.Load(32); uint flags_1 = uint(end_clip.blend != 3u) << uint(1);
param_36.offset = _1368.offset; Alloc _1480;
AnnotatedRef param_37 = out_ref; _1480.offset = _1054.Load(32);
AnnoEndClip param_38 = anno_end_clip; param_38.offset = _1480.offset;
Annotated_EndClip_write(param_36, param_37, param_38); AnnotatedRef param_39 = out_ref;
uint param_40 = flags_1;
AnnoEndClip param_41 = anno_end_clip;
Annotated_EndClip_write(param_38, param_39, param_40, param_41);
} }
} }
if ((tag_word == 9u) || (tag_word == 10u)) if ((tag_word == 9u) || (tag_word == 10u))
@ -709,7 +759,7 @@ void comp_main()
{ {
path_ix = m.path_ix; path_ix = m.path_ix;
} }
_187.Store((clip_out_base + m.clip_ix) * 4 + 8, path_ix); _199.Store((clip_out_base + m.clip_ix) * 4 + 8, path_ix);
} }
out_ref.offset += 40u; out_ref.offset += 40u;
} }

BIN
piet-gpu/shader/gen/kernel4.dxil generated
View file

Binary file not shown.

578
piet-gpu/shader/gen/kernel4.hlsl generated
View file

@ -69,6 +69,16 @@ struct CmdAlpha
float alpha; float alpha;
}; };
struct CmdEndClipRef
{
uint offset;
};
struct CmdEndClip
{
uint blend;
};
struct CmdJumpRef struct CmdJumpRef
{ {
uint offset; uint offset;
@ -132,8 +142,8 @@ struct Config
static const uint3 gl_WorkGroupSize = uint3(8u, 4u, 1u); static const uint3 gl_WorkGroupSize = uint3(8u, 4u, 1u);
RWByteAddressBuffer _202 : register(u0, space0); RWByteAddressBuffer _278 : register(u0, space0);
ByteAddressBuffer _723 : register(t1, space0); ByteAddressBuffer _1521 : register(t1, space0);
RWTexture2D<unorm float4> image_atlas : register(u3, space0); RWTexture2D<unorm float4> image_atlas : register(u3, space0);
RWTexture2D<unorm float4> gradients : register(u4, space0); RWTexture2D<unorm float4> gradients : register(u4, space0);
RWTexture2D<unorm float4> image : register(u2, space0); RWTexture2D<unorm float4> image : register(u2, space0);
@ -160,8 +170,8 @@ float4 spvUnpackUnorm4x8(uint value)
Alloc slice_mem(Alloc a, uint offset, uint size) Alloc slice_mem(Alloc a, uint offset, uint size)
{ {
Alloc _215 = { a.offset + offset }; Alloc _291 = { a.offset + offset };
return _215; return _291;
} }
bool touch_mem(Alloc alloc, uint offset) bool touch_mem(Alloc alloc, uint offset)
@ -177,7 +187,7 @@ uint read_mem(Alloc alloc, uint offset)
{ {
return 0u; return 0u;
} }
uint v = _202.Load(offset * 4 + 8); uint v = _278.Load(offset * 4 + 8);
return v; return v;
} }
@ -186,8 +196,8 @@ CmdTag Cmd_tag(Alloc a, CmdRef ref)
Alloc param = a; Alloc param = a;
uint param_1 = ref.offset >> uint(2); uint param_1 = ref.offset >> uint(2);
uint tag_and_flags = read_mem(param, param_1); uint tag_and_flags = read_mem(param, param_1);
CmdTag _432 = { tag_and_flags & 65535u, tag_and_flags >> uint(16) }; CmdTag _525 = { tag_and_flags & 65535u, tag_and_flags >> uint(16) };
return _432; return _525;
} }
CmdStroke CmdStroke_read(Alloc a, CmdStrokeRef ref) CmdStroke CmdStroke_read(Alloc a, CmdStrokeRef ref)
@ -207,9 +217,9 @@ CmdStroke CmdStroke_read(Alloc a, CmdStrokeRef ref)
CmdStroke Cmd_Stroke_read(Alloc a, CmdRef ref) CmdStroke Cmd_Stroke_read(Alloc a, CmdRef ref)
{ {
CmdStrokeRef _449 = { ref.offset + 4u }; CmdStrokeRef _542 = { ref.offset + 4u };
Alloc param = a; Alloc param = a;
CmdStrokeRef param_1 = _449; CmdStrokeRef param_1 = _542;
return CmdStroke_read(param, param_1); return CmdStroke_read(param, param_1);
} }
@ -245,8 +255,8 @@ TileSeg TileSeg_read(Alloc a, TileSegRef ref)
s.origin = float2(asfloat(raw0), asfloat(raw1)); s.origin = float2(asfloat(raw0), asfloat(raw1));
s._vector = float2(asfloat(raw2), asfloat(raw3)); s._vector = float2(asfloat(raw2), asfloat(raw3));
s.y_edge = asfloat(raw4); s.y_edge = asfloat(raw4);
TileSegRef _572 = { raw5 }; TileSegRef _675 = { raw5 };
s.next = _572; s.next = _675;
return s; return s;
} }
@ -272,9 +282,9 @@ CmdFill CmdFill_read(Alloc a, CmdFillRef ref)
CmdFill Cmd_Fill_read(Alloc a, CmdRef ref) CmdFill Cmd_Fill_read(Alloc a, CmdRef ref)
{ {
CmdFillRef _439 = { ref.offset + 4u }; CmdFillRef _532 = { ref.offset + 4u };
Alloc param = a; Alloc param = a;
CmdFillRef param_1 = _439; CmdFillRef param_1 = _532;
return CmdFill_read(param, param_1); return CmdFill_read(param, param_1);
} }
@ -291,9 +301,9 @@ CmdAlpha CmdAlpha_read(Alloc a, CmdAlphaRef ref)
CmdAlpha Cmd_Alpha_read(Alloc a, CmdRef ref) CmdAlpha Cmd_Alpha_read(Alloc a, CmdRef ref)
{ {
CmdAlphaRef _459 = { ref.offset + 4u }; CmdAlphaRef _552 = { ref.offset + 4u };
Alloc param = a; Alloc param = a;
CmdAlphaRef param_1 = _459; CmdAlphaRef param_1 = _552;
return CmdAlpha_read(param, param_1); return CmdAlpha_read(param, param_1);
} }
@ -310,9 +320,9 @@ CmdColor CmdColor_read(Alloc a, CmdColorRef ref)
CmdColor Cmd_Color_read(Alloc a, CmdRef ref) CmdColor Cmd_Color_read(Alloc a, CmdRef ref)
{ {
CmdColorRef _469 = { ref.offset + 4u }; CmdColorRef _562 = { ref.offset + 4u };
Alloc param = a; Alloc param = a;
CmdColorRef param_1 = _469; CmdColorRef param_1 = _562;
return CmdColor_read(param, param_1); return CmdColor_read(param, param_1);
} }
@ -356,9 +366,9 @@ CmdLinGrad CmdLinGrad_read(Alloc a, CmdLinGradRef ref)
CmdLinGrad Cmd_LinGrad_read(Alloc a, CmdRef ref) CmdLinGrad Cmd_LinGrad_read(Alloc a, CmdRef ref)
{ {
CmdLinGradRef _479 = { ref.offset + 4u }; CmdLinGradRef _572 = { ref.offset + 4u };
Alloc param = a; Alloc param = a;
CmdLinGradRef param_1 = _479; CmdLinGradRef param_1 = _572;
return CmdLinGrad_read(param, param_1); return CmdLinGrad_read(param, param_1);
} }
@ -379,9 +389,9 @@ CmdImage CmdImage_read(Alloc a, CmdImageRef ref)
CmdImage Cmd_Image_read(Alloc a, CmdRef ref) CmdImage Cmd_Image_read(Alloc a, CmdRef ref)
{ {
CmdImageRef _489 = { ref.offset + 4u }; CmdImageRef _582 = { ref.offset + 4u };
Alloc param = a; Alloc param = a;
CmdImageRef param_1 = _489; CmdImageRef param_1 = _582;
return CmdImage_read(param, param_1); return CmdImage_read(param, param_1);
} }
@ -394,10 +404,10 @@ void fillImage(out float4 spvReturnValue[8], uint2 xy, CmdImage cmd_img)
int2 uv = int2(xy + chunk_offset(param)) + cmd_img.offset; int2 uv = int2(xy + chunk_offset(param)) + cmd_img.offset;
float4 fg_rgba = image_atlas[uv]; float4 fg_rgba = image_atlas[uv];
float3 param_1 = fg_rgba.xyz; float3 param_1 = fg_rgba.xyz;
float3 _695 = fromsRGB(param_1); float3 _1493 = fromsRGB(param_1);
fg_rgba.x = _695.x; fg_rgba.x = _1493.x;
fg_rgba.y = _695.y; fg_rgba.y = _1493.y;
fg_rgba.z = _695.z; fg_rgba.z = _1493.z;
rgba[i] = fg_rgba; rgba[i] = fg_rgba;
} }
spvReturnValue = rgba; spvReturnValue = rgba;
@ -418,6 +428,438 @@ uint packsRGB(inout float4 rgba)
return spvPackUnorm4x8(rgba.wzyx); return spvPackUnorm4x8(rgba.wzyx);
} }
CmdEndClip CmdEndClip_read(Alloc a, CmdEndClipRef ref)
{
uint ix = ref.offset >> uint(2);
Alloc param = a;
uint param_1 = ix + 0u;
uint raw0 = read_mem(param, param_1);
CmdEndClip s;
s.blend = raw0;
return s;
}
CmdEndClip Cmd_EndClip_read(Alloc a, CmdRef ref)
{
CmdEndClipRef _592 = { ref.offset + 4u };
Alloc param = a;
CmdEndClipRef param_1 = _592;
return CmdEndClip_read(param, param_1);
}
float3 screen(float3 cb, float3 cs)
{
return (cb + cs) - (cb * cs);
}
float3 hard_light(float3 cb, float3 cs)
{
float3 param = cb;
float3 param_1 = (cs * 2.0f) - 1.0f.xxx;
return lerp(screen(param, param_1), (cb * 2.0f) * cs, float3(bool3(cs.x <= 0.5f.xxx.x, cs.y <= 0.5f.xxx.y, cs.z <= 0.5f.xxx.z)));
}
float color_dodge(float cb, float cs)
{
if (cb == 0.0f)
{
return 0.0f;
}
else
{
if (cs == 1.0f)
{
return 1.0f;
}
else
{
return min(1.0f, cb / (1.0f - cs));
}
}
}
float color_burn(float cb, float cs)
{
if (cb == 1.0f)
{
return 1.0f;
}
else
{
if (cs == 0.0f)
{
return 0.0f;
}
else
{
return 1.0f - min(1.0f, (1.0f - cb) / cs);
}
}
}
float3 soft_light(float3 cb, float3 cs)
{
float3 d = lerp(sqrt(cb), ((((cb * 16.0f) - 12.0f.xxx) * cb) + 4.0f.xxx) * cb, float3(bool3(cb.x <= 0.25f.xxx.x, cb.y <= 0.25f.xxx.y, cb.z <= 0.25f.xxx.z)));
return lerp(cb + (((cs * 2.0f) - 1.0f.xxx) * (d - cb)), cb - (((1.0f.xxx - (cs * 2.0f)) * cb) * (1.0f.xxx - cb)), float3(bool3(cs.x <= 0.5f.xxx.x, cs.y <= 0.5f.xxx.y, cs.z <= 0.5f.xxx.z)));
}
float sat(float3 c)
{
return max(c.x, max(c.y, c.z)) - min(c.x, min(c.y, c.z));
}
void set_sat_inner(inout float cmin, inout float cmid, inout float cmax, float s)
{
if (cmax > cmin)
{
cmid = ((cmid - cmin) * s) / (cmax - cmin);
cmax = s;
}
else
{
cmid = 0.0f;
cmax = 0.0f;
}
cmin = 0.0f;
}
float3 set_sat(inout float3 c, float s)
{
if (c.x <= c.y)
{
if (c.y <= c.z)
{
float param = c.x;
float param_1 = c.y;
float param_2 = c.z;
float param_3 = s;
set_sat_inner(param, param_1, param_2, param_3);
c.x = param;
c.y = param_1;
c.z = param_2;
}
else
{
if (c.x <= c.z)
{
float param_4 = c.x;
float param_5 = c.z;
float param_6 = c.y;
float param_7 = s;
set_sat_inner(param_4, param_5, param_6, param_7);
c.x = param_4;
c.z = param_5;
c.y = param_6;
}
else
{
float param_8 = c.z;
float param_9 = c.x;
float param_10 = c.y;
float param_11 = s;
set_sat_inner(param_8, param_9, param_10, param_11);
c.z = param_8;
c.x = param_9;
c.y = param_10;
}
}
}
else
{
if (c.x <= c.z)
{
float param_12 = c.y;
float param_13 = c.x;
float param_14 = c.z;
float param_15 = s;
set_sat_inner(param_12, param_13, param_14, param_15);
c.y = param_12;
c.x = param_13;
c.z = param_14;
}
else
{
if (c.y <= c.z)
{
float param_16 = c.y;
float param_17 = c.z;
float param_18 = c.x;
float param_19 = s;
set_sat_inner(param_16, param_17, param_18, param_19);
c.y = param_16;
c.z = param_17;
c.x = param_18;
}
else
{
float param_20 = c.z;
float param_21 = c.y;
float param_22 = c.x;
float param_23 = s;
set_sat_inner(param_20, param_21, param_22, param_23);
c.z = param_20;
c.y = param_21;
c.x = param_22;
}
}
}
return c;
}
float lum(float3 c)
{
float3 f = float3(0.300000011920928955078125f, 0.589999973773956298828125f, 0.10999999940395355224609375f);
return dot(c, f);
}
float3 clip_color(inout float3 c)
{
float3 param = c;
float L = lum(param);
float n = min(c.x, min(c.y, c.z));
float x = max(c.x, max(c.y, c.z));
if (n < 0.0f)
{
c = L.xxx + (((c - L.xxx) * L) / (L - n).xxx);
}
if (x > 1.0f)
{
c = L.xxx + (((c - L.xxx) * (1.0f - L)) / (x - L).xxx);
}
return c;
}
float3 set_lum(float3 c, float l)
{
float3 param = c;
float3 param_1 = c + (l - lum(param)).xxx;
float3 _901 = clip_color(param_1);
return _901;
}
float3 mix_blend(float3 cb, float3 cs, uint mode)
{
float3 b = 0.0f.xxx;
switch (mode)
{
case 1u:
{
b = cb * cs;
break;
}
case 2u:
{
float3 param = cb;
float3 param_1 = cs;
b = screen(param, param_1);
break;
}
case 3u:
{
float3 param_2 = cs;
float3 param_3 = cb;
b = hard_light(param_2, param_3);
break;
}
case 4u:
{
b = min(cb, cs);
break;
}
case 5u:
{
b = max(cb, cs);
break;
}
case 6u:
{
float param_4 = cb.x;
float param_5 = cs.x;
float param_6 = cb.y;
float param_7 = cs.y;
float param_8 = cb.z;
float param_9 = cs.z;
b = float3(color_dodge(param_4, param_5), color_dodge(param_6, param_7), color_dodge(param_8, param_9));
break;
}
case 7u:
{
float param_10 = cb.x;
float param_11 = cs.x;
float param_12 = cb.y;
float param_13 = cs.y;
float param_14 = cb.z;
float param_15 = cs.z;
b = float3(color_burn(param_10, param_11), color_burn(param_12, param_13), color_burn(param_14, param_15));
break;
}
case 8u:
{
float3 param_16 = cb;
float3 param_17 = cs;
b = hard_light(param_16, param_17);
break;
}
case 9u:
{
float3 param_18 = cb;
float3 param_19 = cs;
b = soft_light(param_18, param_19);
break;
}
case 10u:
{
b = abs(cb - cs);
break;
}
case 11u:
{
b = (cb + cs) - ((cb * 2.0f) * cs);
break;
}
case 12u:
{
float3 param_20 = cb;
float3 param_21 = cs;
float param_22 = sat(param_20);
float3 _1192 = set_sat(param_21, param_22);
float3 param_23 = cb;
float3 param_24 = _1192;
float param_25 = lum(param_23);
b = set_lum(param_24, param_25);
break;
}
case 13u:
{
float3 param_26 = cs;
float3 param_27 = cb;
float param_28 = sat(param_26);
float3 _1206 = set_sat(param_27, param_28);
float3 param_29 = cb;
float3 param_30 = _1206;
float param_31 = lum(param_29);
b = set_lum(param_30, param_31);
break;
}
case 14u:
{
float3 param_32 = cb;
float3 param_33 = cs;
float param_34 = lum(param_32);
b = set_lum(param_33, param_34);
break;
}
case 15u:
{
float3 param_35 = cs;
float3 param_36 = cb;
float param_37 = lum(param_35);
b = set_lum(param_36, param_37);
break;
}
default:
{
b = cs;
break;
}
}
return b;
}
float4 mix_compose(float3 cb, float3 cs, float ab, float as, uint mode)
{
float fa = 0.0f;
float fb = 0.0f;
switch (mode)
{
case 1u:
{
fa = 1.0f;
fb = 0.0f;
break;
}
case 2u:
{
fa = 0.0f;
fb = 1.0f;
break;
}
case 3u:
{
fa = 1.0f;
fb = 1.0f - as;
break;
}
case 4u:
{
fa = 1.0f - ab;
fb = 1.0f;
break;
}
case 5u:
{
fa = ab;
fb = 0.0f;
break;
}
case 6u:
{
fa = 0.0f;
fb = as;
break;
}
case 7u:
{
fa = 1.0f - ab;
fb = 0.0f;
break;
}
case 8u:
{
fa = 0.0f;
fb = 1.0f - as;
break;
}
case 9u:
{
fa = ab;
fb = 1.0f - as;
break;
}
case 10u:
{
fa = 1.0f - ab;
fb = as;
break;
}
case 11u:
{
fa = 1.0f - ab;
fb = 1.0f - as;
break;
}
case 12u:
{
fa = 1.0f;
fb = 1.0f;
break;
}
case 13u:
{
return float4(max(0.0f.xxxx, ((1.0f.xxxx - (float4(cs, as) * as)) + 1.0f.xxxx) - (float4(cb, ab) * ab)).xyz, max(0.0f, ((1.0f - as) + 1.0f) - ab));
}
case 14u:
{
return float4(min(1.0f.xxxx, (float4(cs, as) * as) + (float4(cb, ab) * ab)).xyz, min(1.0f, as + ab));
}
default:
{
break;
}
}
return (float4(cs, as) * (as * fa)) + (float4(cb, ab) * (ab * fb));
}
CmdJump CmdJump_read(Alloc a, CmdJumpRef ref) CmdJump CmdJump_read(Alloc a, CmdJumpRef ref)
{ {
uint ix = ref.offset >> uint(2); uint ix = ref.offset >> uint(2);
@ -431,24 +873,24 @@ CmdJump CmdJump_read(Alloc a, CmdJumpRef ref)
CmdJump Cmd_Jump_read(Alloc a, CmdRef ref) CmdJump Cmd_Jump_read(Alloc a, CmdRef ref)
{ {
CmdJumpRef _499 = { ref.offset + 4u }; CmdJumpRef _602 = { ref.offset + 4u };
Alloc param = a; Alloc param = a;
CmdJumpRef param_1 = _499; CmdJumpRef param_1 = _602;
return CmdJump_read(param, param_1); return CmdJump_read(param, param_1);
} }
void comp_main() void comp_main()
{ {
uint tile_ix = (gl_WorkGroupID.y * _723.Load(8)) + gl_WorkGroupID.x; uint tile_ix = (gl_WorkGroupID.y * _1521.Load(8)) + gl_WorkGroupID.x;
Alloc _738; Alloc _1536;
_738.offset = _723.Load(24); _1536.offset = _1521.Load(24);
Alloc param; Alloc param;
param.offset = _738.offset; param.offset = _1536.offset;
uint param_1 = tile_ix * 1024u; uint param_1 = tile_ix * 1024u;
uint param_2 = 1024u; uint param_2 = 1024u;
Alloc cmd_alloc = slice_mem(param, param_1, param_2); Alloc cmd_alloc = slice_mem(param, param_1, param_2);
CmdRef _747 = { cmd_alloc.offset }; CmdRef _1545 = { cmd_alloc.offset };
CmdRef cmd_ref = _747; CmdRef cmd_ref = _1545;
uint2 xy_uint = uint2(gl_LocalInvocationID.x + (16u * gl_WorkGroupID.x), gl_LocalInvocationID.y + (16u * gl_WorkGroupID.y)); uint2 xy_uint = uint2(gl_LocalInvocationID.x + (16u * gl_WorkGroupID.x), gl_LocalInvocationID.y + (16u * gl_WorkGroupID.y));
float2 xy = float2(xy_uint); float2 xy = float2(xy_uint);
float4 rgba[8]; float4 rgba[8];
@ -457,7 +899,7 @@ void comp_main()
rgba[i] = 0.0f.xxxx; rgba[i] = 0.0f.xxxx;
} }
uint clip_depth = 0u; uint clip_depth = 0u;
bool mem_ok = _202.Load(4) == 0u; bool mem_ok = _278.Load(4) == 0u;
float df[8]; float df[8];
TileSegRef tile_seg_ref; TileSegRef tile_seg_ref;
float area[8]; float area[8];
@ -482,8 +924,8 @@ void comp_main()
{ {
df[k] = 1000000000.0f; df[k] = 1000000000.0f;
} }
TileSegRef _842 = { stroke.tile_ref }; TileSegRef _1638 = { stroke.tile_ref };
tile_seg_ref = _842; tile_seg_ref = _1638;
do do
{ {
uint param_7 = tile_seg_ref.offset; uint param_7 = tile_seg_ref.offset;
@ -519,8 +961,8 @@ void comp_main()
{ {
area[k_3] = float(fill.backdrop); area[k_3] = float(fill.backdrop);
} }
TileSegRef _964 = { fill.tile_ref }; TileSegRef _1758 = { fill.tile_ref };
tile_seg_ref = _964; tile_seg_ref = _1758;
do do
{ {
uint param_15 = tile_seg_ref.offset; uint param_15 = tile_seg_ref.offset;
@ -609,10 +1051,10 @@ void comp_main()
int x = int(round(clamp(my_d, 0.0f, 1.0f) * 511.0f)); int x = int(round(clamp(my_d, 0.0f, 1.0f) * 511.0f));
float4 fg_rgba = gradients[int2(x, int(lin.index))]; float4 fg_rgba = gradients[int2(x, int(lin.index))];
float3 param_29 = fg_rgba.xyz; float3 param_29 = fg_rgba.xyz;
float3 _1298 = fromsRGB(param_29); float3 _2092 = fromsRGB(param_29);
fg_rgba.x = _1298.x; fg_rgba.x = _2092.x;
fg_rgba.y = _1298.y; fg_rgba.y = _2092.y;
fg_rgba.z = _1298.z; fg_rgba.z = _2092.z;
rgba[k_9] = fg_rgba; rgba[k_9] = fg_rgba;
} }
cmd_ref.offset += 20u; cmd_ref.offset += 20u;
@ -625,9 +1067,9 @@ void comp_main()
CmdImage fill_img = Cmd_Image_read(param_30, param_31); CmdImage fill_img = Cmd_Image_read(param_30, param_31);
uint2 param_32 = xy_uint; uint2 param_32 = xy_uint;
CmdImage param_33 = fill_img; CmdImage param_33 = fill_img;
float4 _1327[8]; float4 _2121[8];
fillImage(_1327, param_32, param_33); fillImage(_2121, param_32, param_33);
float4 img[8] = _1327; float4 img[8] = _2121;
for (uint k_10 = 0u; k_10 < 8u; k_10++) for (uint k_10 = 0u; k_10 < 8u; k_10++)
{ {
float4 fg_k_1 = img[k_10] * area[k_10]; float4 fg_k_1 = img[k_10] * area[k_10];
@ -642,8 +1084,8 @@ void comp_main()
{ {
uint d_2 = min(clip_depth, 127u); uint d_2 = min(clip_depth, 127u);
float4 param_34 = float4(rgba[k_11]); float4 param_34 = float4(rgba[k_11]);
uint _1390 = packsRGB(param_34); uint _2184 = packsRGB(param_34);
blend_stack[d_2][k_11] = _1390; blend_stack[d_2][k_11] = _2184;
rgba[k_11] = 0.0f.xxxx; rgba[k_11] = 0.0f.xxxx;
} }
clip_depth++; clip_depth++;
@ -652,24 +1094,44 @@ void comp_main()
} }
case 9u: case 9u:
{ {
Alloc param_35 = cmd_alloc;
CmdRef param_36 = cmd_ref;
CmdEndClip end_clip = Cmd_EndClip_read(param_35, param_36);
uint blend_mode = end_clip.blend >> uint(8);
uint comp_mode = end_clip.blend & 255u;
clip_depth--; clip_depth--;
for (uint k_12 = 0u; k_12 < 8u; k_12++) for (uint k_12 = 0u; k_12 < 8u; k_12++)
{ {
uint d_3 = min(clip_depth, 127u); uint d_3 = min(clip_depth, 127u);
uint param_35 = blend_stack[d_3][k_12]; uint param_37 = blend_stack[d_3][k_12];
float4 bg = unpacksRGB(param_35); float4 bg = unpacksRGB(param_37);
float4 fg_1 = rgba[k_12] * area[k_12]; float4 fg_1 = rgba[k_12] * area[k_12];
rgba[k_12] = (bg * (1.0f - fg_1.w)) + fg_1; float3 param_38 = bg.xyz;
float3 param_39 = fg_1.xyz;
uint param_40 = blend_mode;
float3 blend = mix_blend(param_38, param_39, param_40);
float4 _2251 = fg_1;
float _2255 = fg_1.w;
float3 _2262 = lerp(_2251.xyz, blend, float((_2255 * bg.w) > 0.0f).xxx);
fg_1.x = _2262.x;
fg_1.y = _2262.y;
fg_1.z = _2262.z;
float3 param_41 = bg.xyz;
float3 param_42 = fg_1.xyz;
float param_43 = bg.w;
float param_44 = fg_1.w;
uint param_45 = comp_mode;
rgba[k_12] = mix_compose(param_41, param_42, param_43, param_44, param_45);
} }
cmd_ref.offset += 4u; cmd_ref.offset += 8u;
break; break;
} }
case 10u: case 10u:
{ {
Alloc param_36 = cmd_alloc; Alloc param_46 = cmd_alloc;
CmdRef param_37 = cmd_ref; CmdRef param_47 = cmd_ref;
CmdRef _1453 = { Cmd_Jump_read(param_36, param_37).new_ref }; CmdRef _2299 = { Cmd_Jump_read(param_46, param_47).new_ref };
cmd_ref = _1453; cmd_ref = _2299;
cmd_alloc.offset = cmd_ref.offset; cmd_alloc.offset = cmd_ref.offset;
break; break;
} }
@ -677,9 +1139,9 @@ void comp_main()
} }
for (uint i_1 = 0u; i_1 < 8u; i_1++) for (uint i_1 = 0u; i_1 < 8u; i_1++)
{ {
uint param_38 = i_1; uint param_48 = i_1;
float3 param_39 = rgba[i_1].xyz; float3 param_49 = rgba[i_1].xyz;
image[int2(xy_uint + chunk_offset(param_38))] = float4(tosRGB(param_39), rgba[i_1].w); image[int2(xy_uint + chunk_offset(param_48))] = float4(tosRGB(param_49), rgba[i_1].w);
} }
} }

6
piet-gpu/shader/gen/kernel4.msl generated
View file

@ -511,7 +511,7 @@ float3 hard_light(thread const float3& cb, thread const float3& cs)
{ {
float3 param = cb; float3 param = cb;
float3 param_1 = (cs * 2.0) - float3(1.0); float3 param_1 = (cs * 2.0) - float3(1.0);
return mix(screen(param, param_1), (cb * 2.0) * cs, select(float3(0.0), float3(1.0), cs <= float3(0.5))); return mix(screen(param, param_1), (cb * 2.0) * cs, float3(cs <= float3(0.5)));
} }
static inline __attribute__((always_inline)) static inline __attribute__((always_inline))
@ -557,8 +557,8 @@ float color_burn(thread const float& cb, thread const float& cs)
static inline __attribute__((always_inline)) static inline __attribute__((always_inline))
float3 soft_light(thread const float3& cb, thread const float3& cs) float3 soft_light(thread const float3& cb, thread const float3& cs)
{ {
float3 d = mix(sqrt(cb), ((((cb * 16.0) - float3(12.0)) * cb) + float3(4.0)) * cb, select(float3(0.0), float3(1.0), cb <= float3(0.25))); float3 d = mix(sqrt(cb), ((((cb * 16.0) - float3(12.0)) * cb) + float3(4.0)) * cb, float3(cb <= float3(0.25)));
return mix(cb + (((cs * 2.0) - float3(1.0)) * (d - cb)), cb - (((float3(1.0) - (cs * 2.0)) * cb) * (float3(1.0) - cb)), select(float3(0.0), float3(1.0), cs <= float3(0.5))); return mix(cb + (((cs * 2.0) - float3(1.0)) * (d - cb)), cb - (((float3(1.0) - (cs * 2.0)) * cb) * (float3(1.0) - cb)), float3(cs <= float3(0.5)));
} }
static inline __attribute__((always_inline)) static inline __attribute__((always_inline))

BIN
piet-gpu/shader/gen/kernel4_gray.dxil generated
View file

Binary file not shown.

576
piet-gpu/shader/gen/kernel4_gray.hlsl generated
View file

@ -69,6 +69,16 @@ struct CmdAlpha
float alpha; float alpha;
}; };
struct CmdEndClipRef
{
uint offset;
};
struct CmdEndClip
{
uint blend;
};
struct CmdJumpRef struct CmdJumpRef
{ {
uint offset; uint offset;
@ -132,8 +142,8 @@ struct Config
static const uint3 gl_WorkGroupSize = uint3(8u, 4u, 1u); static const uint3 gl_WorkGroupSize = uint3(8u, 4u, 1u);
RWByteAddressBuffer _202 : register(u0, space0); RWByteAddressBuffer _278 : register(u0, space0);
ByteAddressBuffer _723 : register(t1, space0); ByteAddressBuffer _1521 : register(t1, space0);
RWTexture2D<unorm float4> image_atlas : register(u3, space0); RWTexture2D<unorm float4> image_atlas : register(u3, space0);
RWTexture2D<unorm float4> gradients : register(u4, space0); RWTexture2D<unorm float4> gradients : register(u4, space0);
RWTexture2D<unorm float> image : register(u2, space0); RWTexture2D<unorm float> image : register(u2, space0);
@ -160,8 +170,8 @@ float4 spvUnpackUnorm4x8(uint value)
Alloc slice_mem(Alloc a, uint offset, uint size) Alloc slice_mem(Alloc a, uint offset, uint size)
{ {
Alloc _215 = { a.offset + offset }; Alloc _291 = { a.offset + offset };
return _215; return _291;
} }
bool touch_mem(Alloc alloc, uint offset) bool touch_mem(Alloc alloc, uint offset)
@ -177,7 +187,7 @@ uint read_mem(Alloc alloc, uint offset)
{ {
return 0u; return 0u;
} }
uint v = _202.Load(offset * 4 + 8); uint v = _278.Load(offset * 4 + 8);
return v; return v;
} }
@ -186,8 +196,8 @@ CmdTag Cmd_tag(Alloc a, CmdRef ref)
Alloc param = a; Alloc param = a;
uint param_1 = ref.offset >> uint(2); uint param_1 = ref.offset >> uint(2);
uint tag_and_flags = read_mem(param, param_1); uint tag_and_flags = read_mem(param, param_1);
CmdTag _432 = { tag_and_flags & 65535u, tag_and_flags >> uint(16) }; CmdTag _525 = { tag_and_flags & 65535u, tag_and_flags >> uint(16) };
return _432; return _525;
} }
CmdStroke CmdStroke_read(Alloc a, CmdStrokeRef ref) CmdStroke CmdStroke_read(Alloc a, CmdStrokeRef ref)
@ -207,9 +217,9 @@ CmdStroke CmdStroke_read(Alloc a, CmdStrokeRef ref)
CmdStroke Cmd_Stroke_read(Alloc a, CmdRef ref) CmdStroke Cmd_Stroke_read(Alloc a, CmdRef ref)
{ {
CmdStrokeRef _449 = { ref.offset + 4u }; CmdStrokeRef _542 = { ref.offset + 4u };
Alloc param = a; Alloc param = a;
CmdStrokeRef param_1 = _449; CmdStrokeRef param_1 = _542;
return CmdStroke_read(param, param_1); return CmdStroke_read(param, param_1);
} }
@ -245,8 +255,8 @@ TileSeg TileSeg_read(Alloc a, TileSegRef ref)
s.origin = float2(asfloat(raw0), asfloat(raw1)); s.origin = float2(asfloat(raw0), asfloat(raw1));
s._vector = float2(asfloat(raw2), asfloat(raw3)); s._vector = float2(asfloat(raw2), asfloat(raw3));
s.y_edge = asfloat(raw4); s.y_edge = asfloat(raw4);
TileSegRef _572 = { raw5 }; TileSegRef _675 = { raw5 };
s.next = _572; s.next = _675;
return s; return s;
} }
@ -272,9 +282,9 @@ CmdFill CmdFill_read(Alloc a, CmdFillRef ref)
CmdFill Cmd_Fill_read(Alloc a, CmdRef ref) CmdFill Cmd_Fill_read(Alloc a, CmdRef ref)
{ {
CmdFillRef _439 = { ref.offset + 4u }; CmdFillRef _532 = { ref.offset + 4u };
Alloc param = a; Alloc param = a;
CmdFillRef param_1 = _439; CmdFillRef param_1 = _532;
return CmdFill_read(param, param_1); return CmdFill_read(param, param_1);
} }
@ -291,9 +301,9 @@ CmdAlpha CmdAlpha_read(Alloc a, CmdAlphaRef ref)
CmdAlpha Cmd_Alpha_read(Alloc a, CmdRef ref) CmdAlpha Cmd_Alpha_read(Alloc a, CmdRef ref)
{ {
CmdAlphaRef _459 = { ref.offset + 4u }; CmdAlphaRef _552 = { ref.offset + 4u };
Alloc param = a; Alloc param = a;
CmdAlphaRef param_1 = _459; CmdAlphaRef param_1 = _552;
return CmdAlpha_read(param, param_1); return CmdAlpha_read(param, param_1);
} }
@ -310,9 +320,9 @@ CmdColor CmdColor_read(Alloc a, CmdColorRef ref)
CmdColor Cmd_Color_read(Alloc a, CmdRef ref) CmdColor Cmd_Color_read(Alloc a, CmdRef ref)
{ {
CmdColorRef _469 = { ref.offset + 4u }; CmdColorRef _562 = { ref.offset + 4u };
Alloc param = a; Alloc param = a;
CmdColorRef param_1 = _469; CmdColorRef param_1 = _562;
return CmdColor_read(param, param_1); return CmdColor_read(param, param_1);
} }
@ -356,9 +366,9 @@ CmdLinGrad CmdLinGrad_read(Alloc a, CmdLinGradRef ref)
CmdLinGrad Cmd_LinGrad_read(Alloc a, CmdRef ref) CmdLinGrad Cmd_LinGrad_read(Alloc a, CmdRef ref)
{ {
CmdLinGradRef _479 = { ref.offset + 4u }; CmdLinGradRef _572 = { ref.offset + 4u };
Alloc param = a; Alloc param = a;
CmdLinGradRef param_1 = _479; CmdLinGradRef param_1 = _572;
return CmdLinGrad_read(param, param_1); return CmdLinGrad_read(param, param_1);
} }
@ -379,9 +389,9 @@ CmdImage CmdImage_read(Alloc a, CmdImageRef ref)
CmdImage Cmd_Image_read(Alloc a, CmdRef ref) CmdImage Cmd_Image_read(Alloc a, CmdRef ref)
{ {
CmdImageRef _489 = { ref.offset + 4u }; CmdImageRef _582 = { ref.offset + 4u };
Alloc param = a; Alloc param = a;
CmdImageRef param_1 = _489; CmdImageRef param_1 = _582;
return CmdImage_read(param, param_1); return CmdImage_read(param, param_1);
} }
@ -394,10 +404,10 @@ void fillImage(out float4 spvReturnValue[8], uint2 xy, CmdImage cmd_img)
int2 uv = int2(xy + chunk_offset(param)) + cmd_img.offset; int2 uv = int2(xy + chunk_offset(param)) + cmd_img.offset;
float4 fg_rgba = image_atlas[uv]; float4 fg_rgba = image_atlas[uv];
float3 param_1 = fg_rgba.xyz; float3 param_1 = fg_rgba.xyz;
float3 _695 = fromsRGB(param_1); float3 _1493 = fromsRGB(param_1);
fg_rgba.x = _695.x; fg_rgba.x = _1493.x;
fg_rgba.y = _695.y; fg_rgba.y = _1493.y;
fg_rgba.z = _695.z; fg_rgba.z = _1493.z;
rgba[i] = fg_rgba; rgba[i] = fg_rgba;
} }
spvReturnValue = rgba; spvReturnValue = rgba;
@ -418,6 +428,438 @@ uint packsRGB(inout float4 rgba)
return spvPackUnorm4x8(rgba.wzyx); return spvPackUnorm4x8(rgba.wzyx);
} }
CmdEndClip CmdEndClip_read(Alloc a, CmdEndClipRef ref)
{
uint ix = ref.offset >> uint(2);
Alloc param = a;
uint param_1 = ix + 0u;
uint raw0 = read_mem(param, param_1);
CmdEndClip s;
s.blend = raw0;
return s;
}
CmdEndClip Cmd_EndClip_read(Alloc a, CmdRef ref)
{
CmdEndClipRef _592 = { ref.offset + 4u };
Alloc param = a;
CmdEndClipRef param_1 = _592;
return CmdEndClip_read(param, param_1);
}
float3 screen(float3 cb, float3 cs)
{
return (cb + cs) - (cb * cs);
}
float3 hard_light(float3 cb, float3 cs)
{
float3 param = cb;
float3 param_1 = (cs * 2.0f) - 1.0f.xxx;
return lerp(screen(param, param_1), (cb * 2.0f) * cs, float3(bool3(cs.x <= 0.5f.xxx.x, cs.y <= 0.5f.xxx.y, cs.z <= 0.5f.xxx.z)));
}
float color_dodge(float cb, float cs)
{
if (cb == 0.0f)
{
return 0.0f;
}
else
{
if (cs == 1.0f)
{
return 1.0f;
}
else
{
return min(1.0f, cb / (1.0f - cs));
}
}
}
float color_burn(float cb, float cs)
{
if (cb == 1.0f)
{
return 1.0f;
}
else
{
if (cs == 0.0f)
{
return 0.0f;
}
else
{
return 1.0f - min(1.0f, (1.0f - cb) / cs);
}
}
}
float3 soft_light(float3 cb, float3 cs)
{
float3 d = lerp(sqrt(cb), ((((cb * 16.0f) - 12.0f.xxx) * cb) + 4.0f.xxx) * cb, float3(bool3(cb.x <= 0.25f.xxx.x, cb.y <= 0.25f.xxx.y, cb.z <= 0.25f.xxx.z)));
return lerp(cb + (((cs * 2.0f) - 1.0f.xxx) * (d - cb)), cb - (((1.0f.xxx - (cs * 2.0f)) * cb) * (1.0f.xxx - cb)), float3(bool3(cs.x <= 0.5f.xxx.x, cs.y <= 0.5f.xxx.y, cs.z <= 0.5f.xxx.z)));
}
float sat(float3 c)
{
return max(c.x, max(c.y, c.z)) - min(c.x, min(c.y, c.z));
}
void set_sat_inner(inout float cmin, inout float cmid, inout float cmax, float s)
{
if (cmax > cmin)
{
cmid = ((cmid - cmin) * s) / (cmax - cmin);
cmax = s;
}
else
{
cmid = 0.0f;
cmax = 0.0f;
}
cmin = 0.0f;
}
float3 set_sat(inout float3 c, float s)
{
if (c.x <= c.y)
{
if (c.y <= c.z)
{
float param = c.x;
float param_1 = c.y;
float param_2 = c.z;
float param_3 = s;
set_sat_inner(param, param_1, param_2, param_3);
c.x = param;
c.y = param_1;
c.z = param_2;
}
else
{
if (c.x <= c.z)
{
float param_4 = c.x;
float param_5 = c.z;
float param_6 = c.y;
float param_7 = s;
set_sat_inner(param_4, param_5, param_6, param_7);
c.x = param_4;
c.z = param_5;
c.y = param_6;
}
else
{
float param_8 = c.z;
float param_9 = c.x;
float param_10 = c.y;
float param_11 = s;
set_sat_inner(param_8, param_9, param_10, param_11);
c.z = param_8;
c.x = param_9;
c.y = param_10;
}
}
}
else
{
if (c.x <= c.z)
{
float param_12 = c.y;
float param_13 = c.x;
float param_14 = c.z;
float param_15 = s;
set_sat_inner(param_12, param_13, param_14, param_15);
c.y = param_12;
c.x = param_13;
c.z = param_14;
}
else
{
if (c.y <= c.z)
{
float param_16 = c.y;
float param_17 = c.z;
float param_18 = c.x;
float param_19 = s;
set_sat_inner(param_16, param_17, param_18, param_19);
c.y = param_16;
c.z = param_17;
c.x = param_18;
}
else
{
float param_20 = c.z;
float param_21 = c.y;
float param_22 = c.x;
float param_23 = s;
set_sat_inner(param_20, param_21, param_22, param_23);
c.z = param_20;
c.y = param_21;
c.x = param_22;
}
}
}
return c;
}
float lum(float3 c)
{
float3 f = float3(0.300000011920928955078125f, 0.589999973773956298828125f, 0.10999999940395355224609375f);
return dot(c, f);
}
float3 clip_color(inout float3 c)
{
float3 param = c;
float L = lum(param);
float n = min(c.x, min(c.y, c.z));
float x = max(c.x, max(c.y, c.z));
if (n < 0.0f)
{
c = L.xxx + (((c - L.xxx) * L) / (L - n).xxx);
}
if (x > 1.0f)
{
c = L.xxx + (((c - L.xxx) * (1.0f - L)) / (x - L).xxx);
}
return c;
}
float3 set_lum(float3 c, float l)
{
float3 param = c;
float3 param_1 = c + (l - lum(param)).xxx;
float3 _901 = clip_color(param_1);
return _901;
}
float3 mix_blend(float3 cb, float3 cs, uint mode)
{
float3 b = 0.0f.xxx;
switch (mode)
{
case 1u:
{
b = cb * cs;
break;
}
case 2u:
{
float3 param = cb;
float3 param_1 = cs;
b = screen(param, param_1);
break;
}
case 3u:
{
float3 param_2 = cs;
float3 param_3 = cb;
b = hard_light(param_2, param_3);
break;
}
case 4u:
{
b = min(cb, cs);
break;
}
case 5u:
{
b = max(cb, cs);
break;
}
case 6u:
{
float param_4 = cb.x;
float param_5 = cs.x;
float param_6 = cb.y;
float param_7 = cs.y;
float param_8 = cb.z;
float param_9 = cs.z;
b = float3(color_dodge(param_4, param_5), color_dodge(param_6, param_7), color_dodge(param_8, param_9));
break;
}
case 7u:
{
float param_10 = cb.x;
float param_11 = cs.x;
float param_12 = cb.y;
float param_13 = cs.y;
float param_14 = cb.z;
float param_15 = cs.z;
b = float3(color_burn(param_10, param_11), color_burn(param_12, param_13), color_burn(param_14, param_15));
break;
}
case 8u:
{
float3 param_16 = cb;
float3 param_17 = cs;
b = hard_light(param_16, param_17);
break;
}
case 9u:
{
float3 param_18 = cb;
float3 param_19 = cs;
b = soft_light(param_18, param_19);
break;
}
case 10u:
{
b = abs(cb - cs);
break;
}
case 11u:
{
b = (cb + cs) - ((cb * 2.0f) * cs);
break;
}
case 12u:
{
float3 param_20 = cb;
float3 param_21 = cs;
float param_22 = sat(param_20);
float3 _1192 = set_sat(param_21, param_22);
float3 param_23 = cb;
float3 param_24 = _1192;
float param_25 = lum(param_23);
b = set_lum(param_24, param_25);
break;
}
case 13u:
{
float3 param_26 = cs;
float3 param_27 = cb;
float param_28 = sat(param_26);
float3 _1206 = set_sat(param_27, param_28);
float3 param_29 = cb;
float3 param_30 = _1206;
float param_31 = lum(param_29);
b = set_lum(param_30, param_31);
break;
}
case 14u:
{
float3 param_32 = cb;
float3 param_33 = cs;
float param_34 = lum(param_32);
b = set_lum(param_33, param_34);
break;
}
case 15u:
{
float3 param_35 = cs;
float3 param_36 = cb;
float param_37 = lum(param_35);
b = set_lum(param_36, param_37);
break;
}
default:
{
b = cs;
break;
}
}
return b;
}
float4 mix_compose(float3 cb, float3 cs, float ab, float as, uint mode)
{
float fa = 0.0f;
float fb = 0.0f;
switch (mode)
{
case 1u:
{
fa = 1.0f;
fb = 0.0f;
break;
}
case 2u:
{
fa = 0.0f;
fb = 1.0f;
break;
}
case 3u:
{
fa = 1.0f;
fb = 1.0f - as;
break;
}
case 4u:
{
fa = 1.0f - ab;
fb = 1.0f;
break;
}
case 5u:
{
fa = ab;
fb = 0.0f;
break;
}
case 6u:
{
fa = 0.0f;
fb = as;
break;
}
case 7u:
{
fa = 1.0f - ab;
fb = 0.0f;
break;
}
case 8u:
{
fa = 0.0f;
fb = 1.0f - as;
break;
}
case 9u:
{
fa = ab;
fb = 1.0f - as;
break;
}
case 10u:
{
fa = 1.0f - ab;
fb = as;
break;
}
case 11u:
{
fa = 1.0f - ab;
fb = 1.0f - as;
break;
}
case 12u:
{
fa = 1.0f;
fb = 1.0f;
break;
}
case 13u:
{
return float4(max(0.0f.xxxx, ((1.0f.xxxx - (float4(cs, as) * as)) + 1.0f.xxxx) - (float4(cb, ab) * ab)).xyz, max(0.0f, ((1.0f - as) + 1.0f) - ab));
}
case 14u:
{
return float4(min(1.0f.xxxx, (float4(cs, as) * as) + (float4(cb, ab) * ab)).xyz, min(1.0f, as + ab));
}
default:
{
break;
}
}
return (float4(cs, as) * (as * fa)) + (float4(cb, ab) * (ab * fb));
}
CmdJump CmdJump_read(Alloc a, CmdJumpRef ref) CmdJump CmdJump_read(Alloc a, CmdJumpRef ref)
{ {
uint ix = ref.offset >> uint(2); uint ix = ref.offset >> uint(2);
@ -431,24 +873,24 @@ CmdJump CmdJump_read(Alloc a, CmdJumpRef ref)
CmdJump Cmd_Jump_read(Alloc a, CmdRef ref) CmdJump Cmd_Jump_read(Alloc a, CmdRef ref)
{ {
CmdJumpRef _499 = { ref.offset + 4u }; CmdJumpRef _602 = { ref.offset + 4u };
Alloc param = a; Alloc param = a;
CmdJumpRef param_1 = _499; CmdJumpRef param_1 = _602;
return CmdJump_read(param, param_1); return CmdJump_read(param, param_1);
} }
void comp_main() void comp_main()
{ {
uint tile_ix = (gl_WorkGroupID.y * _723.Load(8)) + gl_WorkGroupID.x; uint tile_ix = (gl_WorkGroupID.y * _1521.Load(8)) + gl_WorkGroupID.x;
Alloc _738; Alloc _1536;
_738.offset = _723.Load(24); _1536.offset = _1521.Load(24);
Alloc param; Alloc param;
param.offset = _738.offset; param.offset = _1536.offset;
uint param_1 = tile_ix * 1024u; uint param_1 = tile_ix * 1024u;
uint param_2 = 1024u; uint param_2 = 1024u;
Alloc cmd_alloc = slice_mem(param, param_1, param_2); Alloc cmd_alloc = slice_mem(param, param_1, param_2);
CmdRef _747 = { cmd_alloc.offset }; CmdRef _1545 = { cmd_alloc.offset };
CmdRef cmd_ref = _747; CmdRef cmd_ref = _1545;
uint2 xy_uint = uint2(gl_LocalInvocationID.x + (16u * gl_WorkGroupID.x), gl_LocalInvocationID.y + (16u * gl_WorkGroupID.y)); uint2 xy_uint = uint2(gl_LocalInvocationID.x + (16u * gl_WorkGroupID.x), gl_LocalInvocationID.y + (16u * gl_WorkGroupID.y));
float2 xy = float2(xy_uint); float2 xy = float2(xy_uint);
float4 rgba[8]; float4 rgba[8];
@ -457,7 +899,7 @@ void comp_main()
rgba[i] = 0.0f.xxxx; rgba[i] = 0.0f.xxxx;
} }
uint clip_depth = 0u; uint clip_depth = 0u;
bool mem_ok = _202.Load(4) == 0u; bool mem_ok = _278.Load(4) == 0u;
float df[8]; float df[8];
TileSegRef tile_seg_ref; TileSegRef tile_seg_ref;
float area[8]; float area[8];
@ -482,8 +924,8 @@ void comp_main()
{ {
df[k] = 1000000000.0f; df[k] = 1000000000.0f;
} }
TileSegRef _842 = { stroke.tile_ref }; TileSegRef _1638 = { stroke.tile_ref };
tile_seg_ref = _842; tile_seg_ref = _1638;
do do
{ {
uint param_7 = tile_seg_ref.offset; uint param_7 = tile_seg_ref.offset;
@ -519,8 +961,8 @@ void comp_main()
{ {
area[k_3] = float(fill.backdrop); area[k_3] = float(fill.backdrop);
} }
TileSegRef _964 = { fill.tile_ref }; TileSegRef _1758 = { fill.tile_ref };
tile_seg_ref = _964; tile_seg_ref = _1758;
do do
{ {
uint param_15 = tile_seg_ref.offset; uint param_15 = tile_seg_ref.offset;
@ -609,10 +1051,10 @@ void comp_main()
int x = int(round(clamp(my_d, 0.0f, 1.0f) * 511.0f)); int x = int(round(clamp(my_d, 0.0f, 1.0f) * 511.0f));
float4 fg_rgba = gradients[int2(x, int(lin.index))]; float4 fg_rgba = gradients[int2(x, int(lin.index))];
float3 param_29 = fg_rgba.xyz; float3 param_29 = fg_rgba.xyz;
float3 _1298 = fromsRGB(param_29); float3 _2092 = fromsRGB(param_29);
fg_rgba.x = _1298.x; fg_rgba.x = _2092.x;
fg_rgba.y = _1298.y; fg_rgba.y = _2092.y;
fg_rgba.z = _1298.z; fg_rgba.z = _2092.z;
rgba[k_9] = fg_rgba; rgba[k_9] = fg_rgba;
} }
cmd_ref.offset += 20u; cmd_ref.offset += 20u;
@ -625,9 +1067,9 @@ void comp_main()
CmdImage fill_img = Cmd_Image_read(param_30, param_31); CmdImage fill_img = Cmd_Image_read(param_30, param_31);
uint2 param_32 = xy_uint; uint2 param_32 = xy_uint;
CmdImage param_33 = fill_img; CmdImage param_33 = fill_img;
float4 _1327[8]; float4 _2121[8];
fillImage(_1327, param_32, param_33); fillImage(_2121, param_32, param_33);
float4 img[8] = _1327; float4 img[8] = _2121;
for (uint k_10 = 0u; k_10 < 8u; k_10++) for (uint k_10 = 0u; k_10 < 8u; k_10++)
{ {
float4 fg_k_1 = img[k_10] * area[k_10]; float4 fg_k_1 = img[k_10] * area[k_10];
@ -642,8 +1084,8 @@ void comp_main()
{ {
uint d_2 = min(clip_depth, 127u); uint d_2 = min(clip_depth, 127u);
float4 param_34 = float4(rgba[k_11]); float4 param_34 = float4(rgba[k_11]);
uint _1390 = packsRGB(param_34); uint _2184 = packsRGB(param_34);
blend_stack[d_2][k_11] = _1390; blend_stack[d_2][k_11] = _2184;
rgba[k_11] = 0.0f.xxxx; rgba[k_11] = 0.0f.xxxx;
} }
clip_depth++; clip_depth++;
@ -652,24 +1094,44 @@ void comp_main()
} }
case 9u: case 9u:
{ {
Alloc param_35 = cmd_alloc;
CmdRef param_36 = cmd_ref;
CmdEndClip end_clip = Cmd_EndClip_read(param_35, param_36);
uint blend_mode = end_clip.blend >> uint(8);
uint comp_mode = end_clip.blend & 255u;
clip_depth--; clip_depth--;
for (uint k_12 = 0u; k_12 < 8u; k_12++) for (uint k_12 = 0u; k_12 < 8u; k_12++)
{ {
uint d_3 = min(clip_depth, 127u); uint d_3 = min(clip_depth, 127u);
uint param_35 = blend_stack[d_3][k_12]; uint param_37 = blend_stack[d_3][k_12];
float4 bg = unpacksRGB(param_35); float4 bg = unpacksRGB(param_37);
float4 fg_1 = rgba[k_12] * area[k_12]; float4 fg_1 = rgba[k_12] * area[k_12];
rgba[k_12] = (bg * (1.0f - fg_1.w)) + fg_1; float3 param_38 = bg.xyz;
float3 param_39 = fg_1.xyz;
uint param_40 = blend_mode;
float3 blend = mix_blend(param_38, param_39, param_40);
float4 _2251 = fg_1;
float _2255 = fg_1.w;
float3 _2262 = lerp(_2251.xyz, blend, float((_2255 * bg.w) > 0.0f).xxx);
fg_1.x = _2262.x;
fg_1.y = _2262.y;
fg_1.z = _2262.z;
float3 param_41 = bg.xyz;
float3 param_42 = fg_1.xyz;
float param_43 = bg.w;
float param_44 = fg_1.w;
uint param_45 = comp_mode;
rgba[k_12] = mix_compose(param_41, param_42, param_43, param_44, param_45);
} }
cmd_ref.offset += 4u; cmd_ref.offset += 8u;
break; break;
} }
case 10u: case 10u:
{ {
Alloc param_36 = cmd_alloc; Alloc param_46 = cmd_alloc;
CmdRef param_37 = cmd_ref; CmdRef param_47 = cmd_ref;
CmdRef _1453 = { Cmd_Jump_read(param_36, param_37).new_ref }; CmdRef _2299 = { Cmd_Jump_read(param_46, param_47).new_ref };
cmd_ref = _1453; cmd_ref = _2299;
cmd_alloc.offset = cmd_ref.offset; cmd_alloc.offset = cmd_ref.offset;
break; break;
} }
@ -677,8 +1139,8 @@ void comp_main()
} }
for (uint i_1 = 0u; i_1 < 8u; i_1++) for (uint i_1 = 0u; i_1 < 8u; i_1++)
{ {
uint param_38 = i_1; uint param_48 = i_1;
image[int2(xy_uint + chunk_offset(param_38))] = rgba[i_1].w.x; image[int2(xy_uint + chunk_offset(param_48))] = rgba[i_1].w.x;
} }
} }

171
piet-gpu/shader/gen/kernel4_gray.msl generated
View file

@ -454,10 +454,10 @@ spvUnsafeArray<float4, 8> fillImage(thread const uint2& xy, thread const CmdImag
int2 uv = int2(xy + chunk_offset(param)) + cmd_img.offset; int2 uv = int2(xy + chunk_offset(param)) + cmd_img.offset;
float4 fg_rgba = image_atlas.read(uint2(uv)); float4 fg_rgba = image_atlas.read(uint2(uv));
float3 param_1 = fg_rgba.xyz; float3 param_1 = fg_rgba.xyz;
float3 _1495 = fromsRGB(param_1); float3 _1493 = fromsRGB(param_1);
fg_rgba.x = _1495.x; fg_rgba.x = _1493.x;
fg_rgba.y = _1495.y; fg_rgba.y = _1493.y;
fg_rgba.z = _1495.z; fg_rgba.z = _1493.z;
rgba[i] = fg_rgba; rgba[i] = fg_rgba;
} }
return rgba; return rgba;
@ -511,7 +511,7 @@ float3 hard_light(thread const float3& cb, thread const float3& cs)
{ {
float3 param = cb; float3 param = cb;
float3 param_1 = (cs * 2.0) - float3(1.0); float3 param_1 = (cs * 2.0) - float3(1.0);
return mix(screen(param, param_1), (cb * 2.0) * cs, select(float3(0.0), float3(1.0), cs <= float3(0.5))); return mix(screen(param, param_1), (cb * 2.0) * cs, float3(cs <= float3(0.5)));
} }
static inline __attribute__((always_inline)) static inline __attribute__((always_inline))
@ -557,8 +557,8 @@ float color_burn(thread const float& cb, thread const float& cs)
static inline __attribute__((always_inline)) static inline __attribute__((always_inline))
float3 soft_light(thread const float3& cb, thread const float3& cs) float3 soft_light(thread const float3& cb, thread const float3& cs)
{ {
float3 d = mix(sqrt(cb), ((((cb * 16.0) - float3(12.0)) * cb) + float3(4.0)) * cb, select(float3(0.0), float3(1.0), cb <= float3(0.25))); float3 d = mix(sqrt(cb), ((((cb * 16.0) - float3(12.0)) * cb) + float3(4.0)) * cb, float3(cb <= float3(0.25)));
return mix(cb + (((cs * 2.0) - float3(1.0)) * (d - cb)), cb - (((float3(1.0) - (cs * 2.0)) * cb) * (float3(1.0) - cb)), select(float3(0.0), float3(1.0), cs <= float3(0.5))); return mix(cb + (((cs * 2.0) - float3(1.0)) * (d - cb)), cb - (((float3(1.0) - (cs * 2.0)) * cb) * (float3(1.0) - cb)), float3(cs <= float3(0.5)));
} }
static inline __attribute__((always_inline)) static inline __attribute__((always_inline))
@ -568,103 +568,103 @@ float sat(thread const float3& c)
} }
static inline __attribute__((always_inline)) static inline __attribute__((always_inline))
void SetSatInner(thread float& Cmin, thread float& Cmid, thread float& Cmax, thread const float& s) void set_sat_inner(thread float& cmin, thread float& cmid, thread float& cmax, thread const float& s)
{ {
if (Cmax > Cmin) if (cmax > cmin)
{ {
Cmid = ((Cmid - Cmin) * s) / (Cmax - Cmin); cmid = ((cmid - cmin) * s) / (cmax - cmin);
Cmax = s; cmax = s;
} }
else else
{ {
Cmid = 0.0; cmid = 0.0;
Cmax = 0.0; cmax = 0.0;
} }
Cmin = 0.0; cmin = 0.0;
} }
static inline __attribute__((always_inline)) static inline __attribute__((always_inline))
float3 set_sat(thread float3& C, thread const float& s) float3 set_sat(thread float3& c, thread const float& s)
{ {
if (C.x <= C.y) if (c.x <= c.y)
{ {
if (C.y <= C.z) if (c.y <= c.z)
{ {
float param = C.x; float param = c.x;
float param_1 = C.y; float param_1 = c.y;
float param_2 = C.z; float param_2 = c.z;
float param_3 = s; float param_3 = s;
SetSatInner(param, param_1, param_2, param_3); set_sat_inner(param, param_1, param_2, param_3);
C.x = param; c.x = param;
C.y = param_1; c.y = param_1;
C.z = param_2; c.z = param_2;
} }
else else
{ {
if (C.x <= C.z) if (c.x <= c.z)
{ {
float param_4 = C.x; float param_4 = c.x;
float param_5 = C.z; float param_5 = c.z;
float param_6 = C.y; float param_6 = c.y;
float param_7 = s; float param_7 = s;
SetSatInner(param_4, param_5, param_6, param_7); set_sat_inner(param_4, param_5, param_6, param_7);
C.x = param_4; c.x = param_4;
C.z = param_5; c.z = param_5;
C.y = param_6; c.y = param_6;
} }
else else
{ {
float param_8 = C.z; float param_8 = c.z;
float param_9 = C.x; float param_9 = c.x;
float param_10 = C.y; float param_10 = c.y;
float param_11 = s; float param_11 = s;
SetSatInner(param_8, param_9, param_10, param_11); set_sat_inner(param_8, param_9, param_10, param_11);
C.z = param_8; c.z = param_8;
C.x = param_9; c.x = param_9;
C.y = param_10; c.y = param_10;
} }
} }
} }
else else
{ {
if (C.x <= C.z) if (c.x <= c.z)
{ {
float param_12 = C.y; float param_12 = c.y;
float param_13 = C.x; float param_13 = c.x;
float param_14 = C.z; float param_14 = c.z;
float param_15 = s; float param_15 = s;
SetSatInner(param_12, param_13, param_14, param_15); set_sat_inner(param_12, param_13, param_14, param_15);
C.y = param_12; c.y = param_12;
C.x = param_13; c.x = param_13;
C.z = param_14; c.z = param_14;
} }
else else
{ {
if (C.y <= C.z) if (c.y <= c.z)
{ {
float param_16 = C.y; float param_16 = c.y;
float param_17 = C.z; float param_17 = c.z;
float param_18 = C.x; float param_18 = c.x;
float param_19 = s; float param_19 = s;
SetSatInner(param_16, param_17, param_18, param_19); set_sat_inner(param_16, param_17, param_18, param_19);
C.y = param_16; c.y = param_16;
C.z = param_17; c.z = param_17;
C.x = param_18; c.x = param_18;
} }
else else
{ {
float param_20 = C.z; float param_20 = c.z;
float param_21 = C.y; float param_21 = c.y;
float param_22 = C.x; float param_22 = c.x;
float param_23 = s; float param_23 = s;
SetSatInner(param_20, param_21, param_22, param_23); set_sat_inner(param_20, param_21, param_22, param_23);
C.z = param_20; c.z = param_20;
C.y = param_21; c.y = param_21;
C.x = param_22; c.x = param_22;
} }
} }
} }
return C; return c;
} }
static inline __attribute__((always_inline)) static inline __attribute__((always_inline))
@ -696,10 +696,9 @@ static inline __attribute__((always_inline))
float3 set_lum(thread const float3& c, thread const float& l) float3 set_lum(thread const float3& c, thread const float& l)
{ {
float3 param = c; float3 param = c;
float d = l - lum(param); float3 param_1 = c + float3(l - lum(param));
float3 param_1 = c + float3(d); float3 _901 = clip_color(param_1);
float3 _903 = clip_color(param_1); return _901;
return _903;
} }
static inline __attribute__((always_inline)) static inline __attribute__((always_inline))
@ -788,9 +787,9 @@ float3 mix_blend(thread const float3& cb, thread const float3& cs, thread const
float3 param_20 = cb; float3 param_20 = cb;
float3 param_21 = cs; float3 param_21 = cs;
float param_22 = sat(param_20); float param_22 = sat(param_20);
float3 _1194 = set_sat(param_21, param_22); float3 _1192 = set_sat(param_21, param_22);
float3 param_23 = cb; float3 param_23 = cb;
float3 param_24 = _1194; float3 param_24 = _1192;
float param_25 = lum(param_23); float param_25 = lum(param_23);
b = set_lum(param_24, param_25); b = set_lum(param_24, param_25);
break; break;
@ -800,9 +799,9 @@ float3 mix_blend(thread const float3& cb, thread const float3& cs, thread const
float3 param_26 = cs; float3 param_26 = cs;
float3 param_27 = cb; float3 param_27 = cb;
float param_28 = sat(param_26); float param_28 = sat(param_26);
float3 _1208 = set_sat(param_27, param_28); float3 _1206 = set_sat(param_27, param_28);
float3 param_29 = cb; float3 param_29 = cb;
float3 param_30 = _1208; float3 param_30 = _1206;
float param_31 = lum(param_29); float param_31 = lum(param_29);
b = set_lum(param_30, param_31); b = set_lum(param_30, param_31);
break; break;
@ -947,11 +946,11 @@ CmdJump Cmd_Jump_read(thread const Alloc& a, thread const CmdRef& ref, device Me
return CmdJump_read(param, param_1, v_278); return CmdJump_read(param, param_1, v_278);
} }
kernel void main0(device Memory& v_278 [[buffer(0)]], const device ConfigBuf& _1523 [[buffer(1)]], texture2d<float, access::write> image [[texture(2)]], texture2d<float> image_atlas [[texture(3)]], texture2d<float> gradients [[texture(4)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]]) kernel void main0(device Memory& v_278 [[buffer(0)]], const device ConfigBuf& _1521 [[buffer(1)]], texture2d<float, access::write> image [[texture(2)]], texture2d<float> image_atlas [[texture(3)]], texture2d<float> gradients [[texture(4)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
{ {
uint tile_ix = (gl_WorkGroupID.y * _1523.conf.width_in_tiles) + gl_WorkGroupID.x; uint tile_ix = (gl_WorkGroupID.y * _1521.conf.width_in_tiles) + gl_WorkGroupID.x;
Alloc param; Alloc param;
param.offset = _1523.conf.ptcl_alloc.offset; param.offset = _1521.conf.ptcl_alloc.offset;
uint param_1 = tile_ix * 1024u; uint param_1 = tile_ix * 1024u;
uint param_2 = 1024u; uint param_2 = 1024u;
Alloc cmd_alloc = slice_mem(param, param_1, param_2); Alloc cmd_alloc = slice_mem(param, param_1, param_2);
@ -1114,10 +1113,10 @@ kernel void main0(device Memory& v_278 [[buffer(0)]], const device ConfigBuf& _1
int x = int(round(fast::clamp(my_d, 0.0, 1.0) * 511.0)); int x = int(round(fast::clamp(my_d, 0.0, 1.0) * 511.0));
float4 fg_rgba = gradients.read(uint2(int2(x, int(lin.index)))); float4 fg_rgba = gradients.read(uint2(int2(x, int(lin.index))));
float3 param_29 = fg_rgba.xyz; float3 param_29 = fg_rgba.xyz;
float3 _2094 = fromsRGB(param_29); float3 _2092 = fromsRGB(param_29);
fg_rgba.x = _2094.x; fg_rgba.x = _2092.x;
fg_rgba.y = _2094.y; fg_rgba.y = _2092.y;
fg_rgba.z = _2094.z; fg_rgba.z = _2092.z;
rgba[k_9] = fg_rgba; rgba[k_9] = fg_rgba;
} }
cmd_ref.offset += 20u; cmd_ref.offset += 20u;
@ -1146,8 +1145,8 @@ kernel void main0(device Memory& v_278 [[buffer(0)]], const device ConfigBuf& _1
{ {
uint d_2 = min(clip_depth, 127u); uint d_2 = min(clip_depth, 127u);
float4 param_34 = float4(rgba[k_11]); float4 param_34 = float4(rgba[k_11]);
uint _2186 = packsRGB(param_34); uint _2184 = packsRGB(param_34);
blend_stack[d_2][k_11] = _2186; blend_stack[d_2][k_11] = _2184;
rgba[k_11] = float4(0.0); rgba[k_11] = float4(0.0);
} }
clip_depth++; clip_depth++;
@ -1172,12 +1171,12 @@ kernel void main0(device Memory& v_278 [[buffer(0)]], const device ConfigBuf& _1
float3 param_39 = fg_1.xyz; float3 param_39 = fg_1.xyz;
uint param_40 = blend_mode; uint param_40 = blend_mode;
float3 blend = mix_blend(param_38, param_39, param_40); float3 blend = mix_blend(param_38, param_39, param_40);
float4 _2253 = fg_1; float4 _2251 = fg_1;
float _2257 = fg_1.w; float _2255 = fg_1.w;
float3 _2264 = mix(_2253.xyz, blend, float3(float((_2257 * bg.w) > 0.0))); float3 _2262 = mix(_2251.xyz, blend, float3(float((_2255 * bg.w) > 0.0)));
fg_1.x = _2264.x; fg_1.x = _2262.x;
fg_1.y = _2264.y; fg_1.y = _2262.y;
fg_1.z = _2264.z; fg_1.z = _2262.z;
float3 param_41 = bg.xyz; float3 param_41 = bg.xyz;
float3 param_42 = fg_1.xyz; float3 param_42 = fg_1.xyz;
float param_43 = bg.w; float param_43 = bg.w;

BIN
piet-gpu/shader/gen/kernel4_gray.spv generated
View file

Binary file not shown.

BIN
piet-gpu/shader/gen/tile_alloc.dxil generated
View file

Binary file not shown.

77
piet-gpu/shader/gen/tile_alloc.hlsl generated
View file

@ -17,6 +17,7 @@ struct AnnoEndClipRef
struct AnnoEndClip struct AnnoEndClip
{ {
float4 bbox; float4 bbox;
uint blend;
}; };
struct AnnotatedRef struct AnnotatedRef
@ -76,7 +77,7 @@ struct Config
static const uint3 gl_WorkGroupSize = uint3(256u, 1u, 1u); static const uint3 gl_WorkGroupSize = uint3(256u, 1u, 1u);
RWByteAddressBuffer _92 : register(u0, space0); RWByteAddressBuffer _92 : register(u0, space0);
ByteAddressBuffer _305 : register(t1, space0); ByteAddressBuffer _314 : register(t1, space0);
static uint3 gl_LocalInvocationID; static uint3 gl_LocalInvocationID;
static uint3 gl_GlobalInvocationID; static uint3 gl_GlobalInvocationID;
@ -111,8 +112,8 @@ AnnotatedTag Annotated_tag(Alloc a, AnnotatedRef ref)
Alloc param = a; Alloc param = a;
uint param_1 = ref.offset >> uint(2); uint param_1 = ref.offset >> uint(2);
uint tag_and_flags = read_mem(param, param_1); uint tag_and_flags = read_mem(param, param_1);
AnnotatedTag _236 = { tag_and_flags & 65535u, tag_and_flags >> uint(16) }; AnnotatedTag _246 = { tag_and_flags & 65535u, tag_and_flags >> uint(16) };
return _236; return _246;
} }
AnnoEndClip AnnoEndClip_read(Alloc a, AnnoEndClipRef ref) AnnoEndClip AnnoEndClip_read(Alloc a, AnnoEndClipRef ref)
@ -130,16 +131,20 @@ AnnoEndClip AnnoEndClip_read(Alloc a, AnnoEndClipRef ref)
Alloc param_6 = a; Alloc param_6 = a;
uint param_7 = ix + 3u; uint param_7 = ix + 3u;
uint raw3 = read_mem(param_6, param_7); uint raw3 = read_mem(param_6, param_7);
Alloc param_8 = a;
uint param_9 = ix + 4u;
uint raw4 = read_mem(param_8, param_9);
AnnoEndClip s; AnnoEndClip s;
s.bbox = float4(asfloat(raw0), asfloat(raw1), asfloat(raw2), asfloat(raw3)); s.bbox = float4(asfloat(raw0), asfloat(raw1), asfloat(raw2), asfloat(raw3));
s.blend = raw4;
return s; return s;
} }
AnnoEndClip Annotated_EndClip_read(Alloc a, AnnotatedRef ref) AnnoEndClip Annotated_EndClip_read(Alloc a, AnnotatedRef ref)
{ {
AnnoEndClipRef _243 = { ref.offset + 4u }; AnnoEndClipRef _252 = { ref.offset + 4u };
Alloc param = a; Alloc param = a;
AnnoEndClipRef param_1 = _243; AnnoEndClipRef param_1 = _252;
return AnnoEndClip_read(param, param_1); return AnnoEndClip_read(param, param_1);
} }
@ -211,17 +216,17 @@ void comp_main()
{ {
uint th_ix = gl_LocalInvocationID.x; uint th_ix = gl_LocalInvocationID.x;
uint element_ix = gl_GlobalInvocationID.x; uint element_ix = gl_GlobalInvocationID.x;
PathRef _312 = { _305.Load(16) + (element_ix * 12u) }; PathRef _321 = { _314.Load(16) + (element_ix * 12u) };
PathRef path_ref = _312; PathRef path_ref = _321;
AnnotatedRef _321 = { _305.Load(32) + (element_ix * 40u) }; AnnotatedRef _330 = { _314.Load(32) + (element_ix * 40u) };
AnnotatedRef ref = _321; AnnotatedRef ref = _330;
uint tag = 0u; uint tag = 0u;
if (element_ix < _305.Load(0)) if (element_ix < _314.Load(0))
{ {
Alloc _332; Alloc _341;
_332.offset = _305.Load(32); _341.offset = _314.Load(32);
Alloc param; Alloc param;
param.offset = _332.offset; param.offset = _341.offset;
AnnotatedRef param_1 = ref; AnnotatedRef param_1 = ref;
tag = Annotated_tag(param, param_1).tag; tag = Annotated_tag(param, param_1).tag;
} }
@ -237,10 +242,10 @@ void comp_main()
case 4u: case 4u:
case 5u: case 5u:
{ {
Alloc _350; Alloc _359;
_350.offset = _305.Load(32); _359.offset = _314.Load(32);
Alloc param_2; Alloc param_2;
param_2.offset = _350.offset; param_2.offset = _359.offset;
AnnotatedRef param_3 = ref; AnnotatedRef param_3 = ref;
AnnoEndClip clip = Annotated_EndClip_read(param_2, param_3); AnnoEndClip clip = Annotated_EndClip_read(param_2, param_3);
x0 = int(floor(clip.bbox.x * 0.0625f)); x0 = int(floor(clip.bbox.x * 0.0625f));
@ -250,10 +255,10 @@ void comp_main()
break; break;
} }
} }
x0 = clamp(x0, 0, int(_305.Load(8))); x0 = clamp(x0, 0, int(_314.Load(8)));
y0 = clamp(y0, 0, int(_305.Load(12))); y0 = clamp(y0, 0, int(_314.Load(12)));
x1 = clamp(x1, 0, int(_305.Load(8))); x1 = clamp(x1, 0, int(_314.Load(8)));
y1 = clamp(y1, 0, int(_305.Load(12))); y1 = clamp(y1, 0, int(_314.Load(12)));
Path path; Path path;
path.bbox = uint4(uint(x0), uint(y0), uint(x1), uint(y1)); path.bbox = uint4(uint(x0), uint(y0), uint(x1), uint(y1));
uint tile_count = uint((x1 - x0) * (y1 - y0)); uint tile_count = uint((x1 - x0) * (y1 - y0));
@ -276,46 +281,46 @@ void comp_main()
if (th_ix == 255u) if (th_ix == 255u)
{ {
uint param_4 = total_tile_count * 8u; uint param_4 = total_tile_count * 8u;
MallocResult _476 = malloc(param_4); MallocResult _485 = malloc(param_4);
sh_tile_alloc = _476; sh_tile_alloc = _485;
} }
GroupMemoryBarrierWithGroupSync(); GroupMemoryBarrierWithGroupSync();
MallocResult alloc_start = sh_tile_alloc; MallocResult alloc_start = sh_tile_alloc;
bool _487; bool _496;
if (!alloc_start.failed) if (!alloc_start.failed)
{ {
_487 = _92.Load(4) != 0u; _496 = _92.Load(4) != 0u;
} }
else else
{ {
_487 = alloc_start.failed; _496 = alloc_start.failed;
} }
if (_487) if (_496)
{ {
return; return;
} }
if (element_ix < _305.Load(0)) if (element_ix < _314.Load(0))
{ {
uint _500; uint _509;
if (th_ix > 0u) if (th_ix > 0u)
{ {
_500 = sh_tile_count[th_ix - 1u]; _509 = sh_tile_count[th_ix - 1u];
} }
else else
{ {
_500 = 0u; _509 = 0u;
} }
uint tile_subix = _500; uint tile_subix = _509;
Alloc param_5 = alloc_start.alloc; Alloc param_5 = alloc_start.alloc;
uint param_6 = 8u * tile_subix; uint param_6 = 8u * tile_subix;
uint param_7 = 8u * tile_count; uint param_7 = 8u * tile_count;
Alloc tiles_alloc = slice_mem(param_5, param_6, param_7); Alloc tiles_alloc = slice_mem(param_5, param_6, param_7);
TileRef _522 = { tiles_alloc.offset }; TileRef _531 = { tiles_alloc.offset };
path.tiles = _522; path.tiles = _531;
Alloc _527; Alloc _536;
_527.offset = _305.Load(16); _536.offset = _314.Load(16);
Alloc param_8; Alloc param_8;
param_8.offset = _527.offset; param_8.offset = _536.offset;
PathRef param_9 = path_ref; PathRef param_9 = path_ref;
Path param_10 = path; Path param_10 = path;
Path_write(param_8, param_9, param_10); Path_write(param_8, param_9, param_10);