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

Restore gradients and clips

Browse source 
This changes gradients and clips to the new encoding. Lightly tested.
This commit is contained in:
Raph Levien 2021-12-07 17:24:07 -08:00
parent 5585b3563c
commit 49c3a3923b
12 changed files with 147 additions and 117 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

5
piet-gpu/shader/drawtag.h
View file

@ -26,9 +26,10 @@ DrawMonoid map_tag(uint tag_word) {
case Element_FillImage: case Element_FillImage:
return DrawMonoid(1, 0); return DrawMonoid(1, 0);
case Element_BeginClip: case Element_BeginClip:
return DrawMonoid(1, 1); // TODO: endclip should be (0, 1), ie not generate a path. But for now
// we generate a dummy path.
case Element_EndClip: case Element_EndClip:
return DrawMonoid(0, 1); return DrawMonoid(1, 1);
default: default:
return DrawMonoid(0, 0); return DrawMonoid(0, 0);
} }

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

Binary file not shown.

90
piet-gpu/shader/gen/draw_leaf.hlsl
View file

@ -154,14 +154,13 @@ struct Config
static const uint3 gl_WorkGroupSize = uint3(512u, 1u, 1u); static const uint3 gl_WorkGroupSize = uint3(512u, 1u, 1u);
static const DrawMonoid _418 = { 0u, 0u }; static const DrawMonoid _418 = { 0u, 0u };
static const DrawMonoid _443 = { 1u, 0u }; static const DrawMonoid _442 = { 1u, 0u };
static const DrawMonoid _445 = { 1u, 1u }; static const DrawMonoid _444 = { 1u, 1u };
static const DrawMonoid _447 = { 0u, 1u };
RWByteAddressBuffer _201 : register(u0, space0); RWByteAddressBuffer _201 : register(u0, space0);
ByteAddressBuffer _225 : register(t2, space0); ByteAddressBuffer _225 : register(t2, space0);
ByteAddressBuffer _1008 : register(t3, space0); ByteAddressBuffer _1005 : register(t3, space0);
ByteAddressBuffer _1042 : register(t1, space0); ByteAddressBuffer _1039 : register(t1, space0);
static uint3 gl_WorkGroupID; static uint3 gl_WorkGroupID;
static uint3 gl_LocalInvocationID; static uint3 gl_LocalInvocationID;
@ -190,15 +189,12 @@ DrawMonoid map_tag(uint tag_word)
case 5u: case 5u:
case 6u: case 6u:
{ {
return _443; return _442;
} }
case 9u: case 9u:
{
return _445;
}
case 10u: case 10u:
{ {
return _447; return _444;
} }
default: default:
{ {
@ -293,9 +289,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 _808 = { ref.offset + 4u }; AnnoColorRef _805 = { ref.offset + 4u };
Alloc param_3 = a; Alloc param_3 = a;
AnnoColorRef param_4 = _808; AnnoColorRef param_4 = _805;
AnnoColor param_5 = s; AnnoColor param_5 = s;
AnnoColor_write(param_3, param_4, param_5); AnnoColor_write(param_3, param_4, param_5);
} }
@ -369,9 +365,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 _829 = { ref.offset + 4u }; AnnoLinGradientRef _826 = { ref.offset + 4u };
Alloc param_3 = a; Alloc param_3 = a;
AnnoLinGradientRef param_4 = _829; AnnoLinGradientRef param_4 = _826;
AnnoLinGradient param_5 = s; AnnoLinGradient param_5 = s;
AnnoLinGradient_write(param_3, param_4, param_5); AnnoLinGradient_write(param_3, param_4, param_5);
} }
@ -433,9 +429,9 @@ 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 _850 = { ref.offset + 4u }; AnnoImageRef _847 = { ref.offset + 4u };
Alloc param_3 = a; Alloc param_3 = a;
AnnoImageRef param_4 = _850; AnnoImageRef param_4 = _847;
AnnoImage param_5 = s; AnnoImage param_5 = s;
AnnoImage_write(param_3, param_4, param_5); AnnoImage_write(param_3, param_4, param_5);
} }
@ -490,9 +486,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 _871 = { ref.offset + 4u }; AnnoBeginClipRef _868 = { ref.offset + 4u };
Alloc param_3 = a; Alloc param_3 = a;
AnnoBeginClipRef param_4 = _871; AnnoBeginClipRef param_4 = _868;
AnnoBeginClip param_5 = s; AnnoBeginClip param_5 = s;
AnnoBeginClip_write(param_3, param_4, param_5); AnnoBeginClip_write(param_3, param_4, param_5);
} }
@ -531,9 +527,9 @@ void Annotated_EndClip_write(Alloc a, AnnotatedRef ref, AnnoEndClip 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);
AnnoEndClipRef _889 = { ref.offset + 4u }; AnnoEndClipRef _886 = { ref.offset + 4u };
Alloc param_3 = a; Alloc param_3 = a;
AnnoEndClipRef param_4 = _889; AnnoEndClipRef param_4 = _886;
AnnoEndClip param_5 = s; AnnoEndClip param_5 = s;
AnnoEndClip_write(param_3, param_4, param_5); AnnoEndClip_write(param_3, param_4, param_5);
} }
@ -541,8 +537,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 _907 = { ix * 36u }; ElementRef _904 = { ix * 36u };
ElementRef ref = _907; ElementRef ref = _904;
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;
@ -579,11 +575,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 _1014; DrawMonoid _1011;
_1014.path_ix = _1008.Load((gl_WorkGroupID.x - 1u) * 8 + 0); _1011.path_ix = _1005.Load((gl_WorkGroupID.x - 1u) * 8 + 0);
_1014.clip_ix = _1008.Load((gl_WorkGroupID.x - 1u) * 8 + 4); _1011.clip_ix = _1005.Load((gl_WorkGroupID.x - 1u) * 8 + 4);
row.path_ix = _1014.path_ix; row.path_ix = _1011.path_ix;
row.clip_ix = _1014.clip_ix; row.clip_ix = _1011.clip_ix;
} }
if (gl_LocalInvocationID.x > 0u) if (gl_LocalInvocationID.x > 0u)
{ {
@ -592,9 +588,9 @@ 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 = (_1042.Load(44) >> uint(2)) + (out_ix * 2u); uint out_base = (_1039.Load(44) >> uint(2)) + (out_ix * 2u);
AnnotatedRef _1058 = { _1042.Load(32) + (out_ix * 40u) }; AnnotatedRef _1055 = { _1039.Load(32) + (out_ix * 40u) };
AnnotatedRef out_ref = _1058; AnnotatedRef out_ref = _1055;
float4 mat; float4 mat;
float2 translate; float2 translate;
AnnoColor anno_fill; AnnoColor anno_fill;
@ -621,7 +617,7 @@ 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)) if (((tag_word == 4u) || (tag_word == 5u)) || (tag_word == 6u))
{ {
uint bbox_offset = (_1042.Load(40) >> uint(2)) + (6u * (m.path_ix - 1u)); uint bbox_offset = (_1039.Load(40) >> uint(2)) + (6u * (m.path_ix - 1u));
float bbox_l = float(_201.Load(bbox_offset * 4 + 8)) - 32768.0f; float bbox_l = float(_201.Load(bbox_offset * 4 + 8)) - 32768.0f;
float bbox_t = float(_201.Load((bbox_offset + 1u) * 4 + 8)) - 32768.0f; float bbox_t = float(_201.Load((bbox_offset + 1u) * 4 + 8)) - 32768.0f;
float bbox_r = float(_201.Load((bbox_offset + 2u) * 4 + 8)) - 32768.0f; float bbox_r = float(_201.Load((bbox_offset + 2u) * 4 + 8)) - 32768.0f;
@ -632,7 +628,7 @@ void comp_main()
if ((linewidth >= 0.0f) || (tag_word == 5u)) if ((linewidth >= 0.0f) || (tag_word == 5u))
{ {
uint trans_ix = _201.Load((bbox_offset + 5u) * 4 + 8); uint trans_ix = _201.Load((bbox_offset + 5u) * 4 + 8);
uint t = (_1042.Load(36) >> uint(2)) + (6u * trans_ix); uint t = (_1039.Load(36) >> uint(2)) + (6u * trans_ix);
mat = asfloat(uint4(_201.Load(t * 4 + 8), _201.Load((t + 1u) * 4 + 8), _201.Load((t + 2u) * 4 + 8), _201.Load((t + 3u) * 4 + 8))); mat = asfloat(uint4(_201.Load(t * 4 + 8), _201.Load((t + 1u) * 4 + 8), _201.Load((t + 2u) * 4 + 8), _201.Load((t + 3u) * 4 + 8)));
if (tag_word == 5u) if (tag_word == 5u)
{ {
@ -653,9 +649,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 _1261; Alloc _1258;
_1261.offset = _1042.Load(32); _1258.offset = _1039.Load(32);
param_18.offset = _1261.offset; param_18.offset = _1258.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;
@ -678,9 +674,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 _1357; Alloc _1354;
_1357.offset = _1042.Load(32); _1354.offset = _1039.Load(32);
param_23.offset = _1357.offset; param_23.offset = _1354.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;
@ -695,9 +691,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 _1385; Alloc _1382;
_1385.offset = _1042.Load(32); _1382.offset = _1039.Load(32);
param_28.offset = _1385.offset; param_28.offset = _1382.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;
@ -714,9 +710,9 @@ void comp_main()
Clip begin_clip = Element_BeginClip_read(param_32); Clip begin_clip = Element_BeginClip_read(param_32);
anno_begin_clip.bbox = begin_clip.bbox; anno_begin_clip.bbox = begin_clip.bbox;
anno_begin_clip.linewidth = 0.0f; anno_begin_clip.linewidth = 0.0f;
Alloc _1413; Alloc _1410;
_1413.offset = _1042.Load(32); _1410.offset = _1039.Load(32);
param_33.offset = _1413.offset; param_33.offset = _1410.offset;
AnnotatedRef param_34 = out_ref; AnnotatedRef param_34 = out_ref;
uint param_35 = 0u; uint param_35 = 0u;
AnnoBeginClip param_36 = anno_begin_clip; AnnoBeginClip param_36 = anno_begin_clip;
@ -729,9 +725,9 @@ void comp_main()
ElementRef param_37 = this_ref; ElementRef param_37 = this_ref;
Clip end_clip = Element_EndClip_read(param_37); Clip end_clip = Element_EndClip_read(param_37);
anno_end_clip.bbox = end_clip.bbox; anno_end_clip.bbox = end_clip.bbox;
Alloc _1438; Alloc _1435;
_1438.offset = _1042.Load(32); _1435.offset = _1039.Load(32);
param_38.offset = _1438.offset; param_38.offset = _1435.offset;
AnnotatedRef param_39 = out_ref; AnnotatedRef param_39 = out_ref;
AnnoEndClip param_40 = anno_end_clip; AnnoEndClip param_40 = anno_end_clip;
Annotated_EndClip_write(param_38, param_39, param_40); Annotated_EndClip_write(param_38, param_39, param_40);

31
piet-gpu/shader/gen/draw_leaf.msl
View file

@ -251,12 +251,9 @@ DrawMonoid map_tag(thread const uint& tag_word)
return DrawMonoid{ 1u, 0u }; return DrawMonoid{ 1u, 0u };
} }
case 9u: case 9u:
{
return DrawMonoid{ 1u, 1u };
}
case 10u: case 10u:
{ {
return DrawMonoid{ 0u, 1u }; return DrawMonoid{ 1u, 1u };
} }
default: default:
{ {
@ -609,7 +606,7 @@ void Annotated_EndClip_write(thread const Alloc& a, thread const AnnotatedRef& r
AnnoEndClip_write(param_3, param_4, param_5, v_201); AnnoEndClip_write(param_3, param_4, param_5, v_201);
} }
kernel void main0(device Memory& v_201 [[buffer(0)]], const device ConfigBuf& _1042 [[buffer(1)]], const device SceneBuf& v_225 [[buffer(2)]], const device ParentBuf& _1008 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]]) kernel void main0(device Memory& v_201 [[buffer(0)]], const device ConfigBuf& _1039 [[buffer(1)]], const device SceneBuf& v_225 [[buffer(2)]], const device ParentBuf& _1005 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
{ {
threadgroup DrawMonoid sh_scratch[512]; threadgroup DrawMonoid sh_scratch[512];
uint ix = gl_GlobalInvocationID.x * 8u; uint ix = gl_GlobalInvocationID.x * 8u;
@ -650,9 +647,9 @@ kernel void main0(device Memory& v_201 [[buffer(0)]], const device ConfigBuf& _1
DrawMonoid row = tag_monoid_identity(); DrawMonoid row = tag_monoid_identity();
if (gl_WorkGroupID.x > 0u) if (gl_WorkGroupID.x > 0u)
{ {
uint _1011 = gl_WorkGroupID.x - 1u; uint _1008 = gl_WorkGroupID.x - 1u;
row.path_ix = _1008.parent[_1011].path_ix; row.path_ix = _1005.parent[_1008].path_ix;
row.clip_ix = _1008.parent[_1011].clip_ix; row.clip_ix = _1005.parent[_1008].clip_ix;
} }
if (gl_LocalInvocationID.x > 0u) if (gl_LocalInvocationID.x > 0u)
{ {
@ -661,8 +658,8 @@ kernel void main0(device Memory& v_201 [[buffer(0)]], const device ConfigBuf& _1
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 = (_1042.conf.drawmonoid_alloc.offset >> uint(2)) + (out_ix * 2u); uint out_base = (_1039.conf.drawmonoid_alloc.offset >> uint(2)) + (out_ix * 2u);
AnnotatedRef out_ref = AnnotatedRef{ _1042.conf.anno_alloc.offset + (out_ix * 40u) }; AnnotatedRef out_ref = AnnotatedRef{ _1039.conf.anno_alloc.offset + (out_ix * 40u) };
float4 mat; float4 mat;
float2 translate; float2 translate;
AnnoColor anno_fill; AnnoColor anno_fill;
@ -689,7 +686,7 @@ kernel void main0(device Memory& v_201 [[buffer(0)]], const device ConfigBuf& _1
tag_word = Element_tag(param_16, v_225).tag; tag_word = Element_tag(param_16, v_225).tag;
if (((tag_word == 4u) || (tag_word == 5u)) || (tag_word == 6u)) if (((tag_word == 4u) || (tag_word == 5u)) || (tag_word == 6u))
{ {
uint bbox_offset = (_1042.conf.bbox_alloc.offset >> uint(2)) + (6u * (m.path_ix - 1u)); uint bbox_offset = (_1039.conf.bbox_alloc.offset >> uint(2)) + (6u * (m.path_ix - 1u));
float bbox_l = float(v_201.memory[bbox_offset]) - 32768.0; float bbox_l = float(v_201.memory[bbox_offset]) - 32768.0;
float bbox_t = float(v_201.memory[bbox_offset + 1u]) - 32768.0; float bbox_t = float(v_201.memory[bbox_offset + 1u]) - 32768.0;
float bbox_r = float(v_201.memory[bbox_offset + 2u]) - 32768.0; float bbox_r = float(v_201.memory[bbox_offset + 2u]) - 32768.0;
@ -700,7 +697,7 @@ kernel void main0(device Memory& v_201 [[buffer(0)]], const device ConfigBuf& _1
if ((linewidth >= 0.0) || (tag_word == 5u)) if ((linewidth >= 0.0) || (tag_word == 5u))
{ {
uint trans_ix = v_201.memory[bbox_offset + 5u]; uint trans_ix = v_201.memory[bbox_offset + 5u];
uint t = (_1042.conf.trans_alloc.offset >> uint(2)) + (6u * trans_ix); uint t = (_1039.conf.trans_alloc.offset >> uint(2)) + (6u * trans_ix);
mat = as_type<float4>(uint4(v_201.memory[t], v_201.memory[t + 1u], v_201.memory[t + 2u], v_201.memory[t + 3u])); mat = as_type<float4>(uint4(v_201.memory[t], v_201.memory[t + 1u], v_201.memory[t + 2u], v_201.memory[t + 3u]));
if (tag_word == 5u) if (tag_word == 5u)
{ {
@ -721,7 +718,7 @@ kernel void main0(device Memory& v_201 [[buffer(0)]], const device ConfigBuf& _1
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;
param_18.offset = _1042.conf.anno_alloc.offset; param_18.offset = _1039.conf.anno_alloc.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;
@ -744,7 +741,7 @@ kernel void main0(device Memory& v_201 [[buffer(0)]], const device ConfigBuf& _1
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));
param_23.offset = _1042.conf.anno_alloc.offset; param_23.offset = _1039.conf.anno_alloc.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;
@ -759,7 +756,7 @@ kernel void main0(device Memory& v_201 [[buffer(0)]], const device ConfigBuf& _1
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;
param_28.offset = _1042.conf.anno_alloc.offset; param_28.offset = _1039.conf.anno_alloc.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;
@ -776,7 +773,7 @@ kernel void main0(device Memory& v_201 [[buffer(0)]], const device ConfigBuf& _1
Clip begin_clip = Element_BeginClip_read(param_32, v_225); Clip begin_clip = Element_BeginClip_read(param_32, v_225);
anno_begin_clip.bbox = begin_clip.bbox; anno_begin_clip.bbox = begin_clip.bbox;
anno_begin_clip.linewidth = 0.0; anno_begin_clip.linewidth = 0.0;
param_33.offset = _1042.conf.anno_alloc.offset; param_33.offset = _1039.conf.anno_alloc.offset;
AnnotatedRef param_34 = out_ref; AnnotatedRef param_34 = out_ref;
uint param_35 = 0u; uint param_35 = 0u;
AnnoBeginClip param_36 = anno_begin_clip; AnnoBeginClip param_36 = anno_begin_clip;
@ -789,7 +786,7 @@ kernel void main0(device Memory& v_201 [[buffer(0)]], const device ConfigBuf& _1
ElementRef param_37 = this_ref; ElementRef param_37 = this_ref;
Clip end_clip = Element_EndClip_read(param_37, v_225); Clip end_clip = Element_EndClip_read(param_37, v_225);
anno_end_clip.bbox = end_clip.bbox; anno_end_clip.bbox = end_clip.bbox;
param_38.offset = _1042.conf.anno_alloc.offset; param_38.offset = _1039.conf.anno_alloc.offset;
AnnotatedRef param_39 = out_ref; AnnotatedRef param_39 = out_ref;
AnnoEndClip param_40 = anno_end_clip; AnnoEndClip param_40 = anno_end_clip;
Annotated_EndClip_write(param_38, param_39, param_40, v_201); Annotated_EndClip_write(param_38, param_39, param_40, v_201);

BIN
piet-gpu/shader/gen/draw_leaf.spv
View file

Binary file not shown.

BIN
piet-gpu/shader/gen/draw_reduce.dxil
View file

Binary file not shown.

30
piet-gpu/shader/gen/draw_reduce.hlsl
View file

@ -44,15 +44,14 @@ struct Config
uint pathseg_offset; uint pathseg_offset;
}; };
static const DrawMonoid _88 = { 1u, 0u }; static const DrawMonoid _87 = { 1u, 0u };
static const DrawMonoid _90 = { 1u, 1u }; static const DrawMonoid _89 = { 1u, 1u };
static const DrawMonoid _92 = { 0u, 1u }; static const DrawMonoid _91 = { 0u, 0u };
static const DrawMonoid _94 = { 0u, 0u };
ByteAddressBuffer _46 : register(t2, space0); ByteAddressBuffer _46 : register(t2, space0);
RWByteAddressBuffer _203 : register(u3, space0); RWByteAddressBuffer _200 : register(u3, space0);
RWByteAddressBuffer _217 : register(u0, space0); RWByteAddressBuffer _214 : register(u0, space0);
ByteAddressBuffer _223 : register(t1, space0); ByteAddressBuffer _220 : register(t1, space0);
static uint3 gl_WorkGroupID; static uint3 gl_WorkGroupID;
static uint3 gl_LocalInvocationID; static uint3 gl_LocalInvocationID;
@ -81,19 +80,16 @@ DrawMonoid map_tag(uint tag_word)
case 5u: case 5u:
case 6u: case 6u:
{ {
return _88; return _87;
} }
case 9u: case 9u:
{
return _90;
}
case 10u: case 10u:
{ {
return _92; return _89;
} }
default: default:
{ {
return _94; return _91;
} }
} }
} }
@ -115,8 +111,8 @@ DrawMonoid combine_tag_monoid(DrawMonoid a, DrawMonoid b)
void comp_main() void comp_main()
{ {
uint ix = gl_GlobalInvocationID.x * 8u; uint ix = gl_GlobalInvocationID.x * 8u;
ElementRef _110 = { ix * 36u }; ElementRef _107 = { ix * 36u };
ElementRef ref = _110; ElementRef ref = _107;
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;
@ -148,8 +144,8 @@ void comp_main()
} }
if (gl_LocalInvocationID.x == 0u) if (gl_LocalInvocationID.x == 0u)
{ {
_203.Store(gl_WorkGroupID.x * 8 + 0, agg.path_ix); _200.Store(gl_WorkGroupID.x * 8 + 0, agg.path_ix);
_203.Store(gl_WorkGroupID.x * 8 + 4, agg.clip_ix); _200.Store(gl_WorkGroupID.x * 8 + 4, agg.clip_ix);
} }
} }

11
piet-gpu/shader/gen/draw_reduce.msl
View file

@ -98,12 +98,9 @@ DrawMonoid map_tag(thread const uint& tag_word)
return DrawMonoid{ 1u, 0u }; return DrawMonoid{ 1u, 0u };
} }
case 9u: case 9u:
{
return DrawMonoid{ 1u, 1u };
}
case 10u: case 10u:
{ {
return DrawMonoid{ 0u, 1u }; return DrawMonoid{ 1u, 1u };
} }
default: default:
{ {
@ -127,7 +124,7 @@ DrawMonoid combine_tag_monoid(thread const DrawMonoid& a, thread const DrawMonoi
return c; return c;
} }
kernel void main0(const device SceneBuf& v_46 [[buffer(2)]], device OutBuf& _203 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]]) kernel void main0(const device SceneBuf& v_46 [[buffer(2)]], device OutBuf& _200 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
{ {
threadgroup DrawMonoid sh_scratch[512]; threadgroup DrawMonoid sh_scratch[512];
uint ix = gl_GlobalInvocationID.x * 8u; uint ix = gl_GlobalInvocationID.x * 8u;
@ -163,8 +160,8 @@ kernel void main0(const device SceneBuf& v_46 [[buffer(2)]], device OutBuf& _203
} }
if (gl_LocalInvocationID.x == 0u) if (gl_LocalInvocationID.x == 0u)
{ {
_203.outbuf[gl_WorkGroupID.x].path_ix = agg.path_ix; _200.outbuf[gl_WorkGroupID.x].path_ix = agg.path_ix;
_203.outbuf[gl_WorkGroupID.x].clip_ix = agg.clip_ix; _200.outbuf[gl_WorkGroupID.x].clip_ix = agg.clip_ix;
} }
} }

BIN
piet-gpu/shader/gen/draw_reduce.spv
View file

Binary file not shown.

60
piet-gpu/src/encoder.rs
View file

@ -62,6 +62,9 @@ const DRAWOBJ_PART_SIZE: usize = 4096;
// Element struct in piet-gpu-types; that's pretty much going away. // Element struct in piet-gpu-types; that's pretty much going away.
const ELEMENT_FILLCOLOR: u32 = 4; const ELEMENT_FILLCOLOR: u32 = 4;
const ELEMENT_FILLLINGRADIENT: u32 = 5;
const ELEMENT_BEGINCLIP: u32 = 9;
const ELEMENT_ENDCLIP: u32 = 10;
#[repr(C)] #[repr(C)]
#[derive(Clone, Copy, Debug, Default, Zeroable, Pod)] #[derive(Clone, Copy, Debug, Default, Zeroable, Pod)]
@ -71,6 +74,24 @@ pub struct FillColor {
padding: [u32; 7], padding: [u32; 7],
} }
#[repr(C)]
#[derive(Clone, Copy, Debug, Default, Zeroable, Pod)]
pub struct FillLinGradient {
tag: u32,
index: u32,
p0: [f32; 2],
p1: [f32; 2],
padding: [u32; 3],
}
#[repr(C)]
#[derive(Clone, Copy, Debug, Default, Zeroable, Pod)]
pub struct Clip {
tag: u32,
bbox: [f32; 4],
padding: [u32; 4],
}
impl Encoder { impl Encoder {
pub fn new() -> Encoder { pub fn new() -> Encoder {
Encoder { Encoder {
@ -116,6 +137,45 @@ impl Encoder {
self.drawobj_stream.extend(bytemuck::bytes_of(&element)); self.drawobj_stream.extend(bytemuck::bytes_of(&element));
} }
/// Encode a fill linear gradient draw object.
///
/// This should be encoded after a path.
pub fn fill_lin_gradient(&mut self, index: u32, p0: [f32; 2], p1: [f32; 2]) {
let element = FillLinGradient {
tag: ELEMENT_FILLLINGRADIENT,
index,
p0,
p1,
..Default::default()
};
self.drawobj_stream.extend(bytemuck::bytes_of(&element));
}
/// Start a clip and return a save point to be filled in later.
pub fn begin_clip(&mut self) -> usize {
let saved = self.drawobj_stream.len();
let element = Clip {
tag: ELEMENT_BEGINCLIP,
..Default::default()
};
self.drawobj_stream.extend(bytemuck::bytes_of(&element));
saved
}
pub fn end_clip(&mut self, bbox: [f32; 4], save_point: usize) {
let element = Clip {
tag: ELEMENT_ENDCLIP,
bbox,
..Default::default()
};
self.drawobj_stream[save_point + 4..save_point + 20]
.clone_from_slice(bytemuck::bytes_of(&bbox));
self.drawobj_stream.extend(bytemuck::bytes_of(&element));
// This is a dummy path, and will go away with the new clip impl.
self.tag_stream.push(0x10);
self.n_path += 1;
}
/// Return a config for the element processing pipeline. /// Return a config for the element processing pipeline.
/// ///
/// This does not include further pipeline processing. Also returns the /// This does not include further pipeline processing. Also returns the

34
piet-gpu/src/render_ctx.rs
View file

@ -11,7 +11,7 @@ use piet::{
use piet_gpu_hal::BufWrite; use piet_gpu_hal::BufWrite;
use piet_gpu_types::encoder::{Encode, Encoder}; use piet_gpu_types::encoder::{Encode, Encoder};
use piet_gpu_types::scene::{Clip, Element, FillLinGradient, SetFillMode}; use piet_gpu_types::scene::{Element, SetFillMode};
use crate::gradient::{LinearGradient, RampCache}; use crate::gradient::{LinearGradient, RampCache};
use crate::text::Font; use crate::text::Font;
@ -64,8 +64,8 @@ struct State {
} }
struct ClipElement { struct ClipElement {
/// Index of BeginClip element in element vec, for bbox fixup. /// Byte offset of BeginClip element in element vec, for bbox fixup.
begin_ix: usize, save_point: usize,
bbox: Option<Rect>, bbox: Option<Rect>,
} }
@ -242,21 +242,17 @@ impl RenderContext for PietGpuRenderContext {
fn fill_even_odd(&mut self, _shape: impl Shape, _brush: &impl IntoBrush<Self>) {} fn fill_even_odd(&mut self, _shape: impl Shape, _brush: &impl IntoBrush<Self>) {}
fn clip(&mut self, shape: impl Shape) { fn clip(&mut self, shape: impl Shape) {
self.set_fill_mode(FillMode::Nonzero); self.encode_linewidth(-1.0);
let path = shape.path_elements(TOLERANCE); let path = shape.path_elements(TOLERANCE);
self.encode_path(path, true); self.encode_path(path, true);
let begin_ix = self.elements.len(); let save_point = self.new_encoder.begin_clip();
self.elements.push(Element::BeginClip(Clip {
bbox: Default::default(),
}));
if self.clip_stack.len() >= MAX_BLEND_STACK { if self.clip_stack.len() >= MAX_BLEND_STACK {
panic!("Maximum clip/blend stack size {} exceeded", MAX_BLEND_STACK); panic!("Maximum clip/blend stack size {} exceeded", MAX_BLEND_STACK);
} }
self.clip_stack.push(ClipElement { self.clip_stack.push(ClipElement {
bbox: None, bbox: None,
begin_ix, save_point,
}); });
self.path_count += 1;
if let Some(tos) = self.state_stack.last_mut() { if let Some(tos) = self.state_stack.last_mut() {
tos.n_clip += 1; tos.n_clip += 1;
} }
@ -405,14 +401,7 @@ impl PietGpuRenderContext {
let tos = self.clip_stack.pop().unwrap(); let tos = self.clip_stack.pop().unwrap();
let bbox = tos.bbox.unwrap_or_default(); let bbox = tos.bbox.unwrap_or_default();
let bbox_f32_4 = rect_to_f32_4(bbox); let bbox_f32_4 = rect_to_f32_4(bbox);
self.elements self.new_encoder.end_clip(bbox_f32_4, tos.save_point);
.push(Element::EndClip(Clip { bbox: bbox_f32_4 }));
self.path_count += 1;
if let Element::BeginClip(begin_clip) = &mut self.elements[tos.begin_ix] {
begin_clip.bbox = bbox_f32_4;
} else {
unreachable!("expected BeginClip, not found");
}
if let Some(bbox) = tos.bbox { if let Some(bbox) = tos.bbox {
self.union_bbox(bbox); self.union_bbox(bbox);
} }
@ -468,13 +457,8 @@ impl PietGpuRenderContext {
self.new_encoder.fill_color(*rgba_color); self.new_encoder.fill_color(*rgba_color);
} }
PietGpuBrush::LinGradient(lin) => { PietGpuBrush::LinGradient(lin) => {
let fill_lin = FillLinGradient { self.new_encoder
index: lin.ramp_id, .fill_lin_gradient(lin.ramp_id, lin.start, lin.end);
p0: lin.start,
p1: lin.end,
};
self.elements.push(Element::FillLinGradient(fill_lin));
self.path_count += 1;
} }
} }
} }

3
piet-gpu/src/text.rs
View file

@ -12,8 +12,8 @@ use piet::{
use crate::encoder::GlyphEncoder; use crate::encoder::GlyphEncoder;
use crate::render_ctx::{self, FillMode}; use crate::render_ctx::{self, FillMode};
use crate::PietGpuRenderContext;
use crate::stages::Transform; use crate::stages::Transform;
use crate::PietGpuRenderContext;
// This is very much a hack to get things working. // This is very much a hack to get things working.
// On Windows, can set this to "c:\\Windows\\Fonts\\seguiemj.ttf" to get color emoji // On Windows, can set this to "c:\\Windows\\Fonts\\seguiemj.ttf" to get color emoji
@ -51,7 +51,6 @@ struct Glyph {
y: f32, y: f32,
} }
struct TextRenderCtx<'a> { struct TextRenderCtx<'a> {
scaler: Scaler<'a>, scaler: Scaler<'a>,
} }