From 49c3a3923b54c2aad224f52f6c648d0492642c93 Mon Sep 17 00:00:00 2001 From: Raph Levien Date: Tue, 7 Dec 2021 17:24:07 -0800 Subject: [PATCH] Restore gradients and clips This changes gradients and clips to the new encoding. Lightly tested. --- piet-gpu/shader/drawtag.h | 5 +- piet-gpu/shader/gen/draw_leaf.dxil | Bin 6908 -> 6880 bytes piet-gpu/shader/gen/draw_leaf.hlsl | 90 +++++++++++++-------------- piet-gpu/shader/gen/draw_leaf.msl | 31 +++++---- piet-gpu/shader/gen/draw_leaf.spv | Bin 38596 -> 38560 bytes piet-gpu/shader/gen/draw_reduce.dxil | Bin 3956 -> 3940 bytes piet-gpu/shader/gen/draw_reduce.hlsl | 30 ++++----- piet-gpu/shader/gen/draw_reduce.msl | 11 ++-- piet-gpu/shader/gen/draw_reduce.spv | Bin 6908 -> 6872 bytes piet-gpu/src/encoder.rs | 60 ++++++++++++++++++ piet-gpu/src/render_ctx.rs | 34 +++------- piet-gpu/src/text.rs | 3 +- 12 files changed, 147 insertions(+), 117 deletions(-) diff --git a/piet-gpu/shader/drawtag.h b/piet-gpu/shader/drawtag.h index a9e8a1d..17105f6 100644 --- a/piet-gpu/shader/drawtag.h +++ b/piet-gpu/shader/drawtag.h @@ -26,9 +26,10 @@ DrawMonoid map_tag(uint tag_word) { case Element_FillImage: return DrawMonoid(1, 0); 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: - return DrawMonoid(0, 1); + return DrawMonoid(1, 1); default: return DrawMonoid(0, 0); } diff --git a/piet-gpu/shader/gen/draw_leaf.dxil b/piet-gpu/shader/gen/draw_leaf.dxil index f95e5bcfcdd97577ea4addb88b9a791740cf430c..86b37e940da57e098366edaadf2139a3bc786b49 100644 GIT binary patch delta 2528 zcmb_eeNYqW8Gkp)E=fqtCJ+Krxa0r@3^qYwh=;HQ!l8l<90?yLutZTnt`|Pe;%9b) zC@54gMpE?70LL_UT8CEehITsF1j6Z|okOY)w$uqKH_>CJjz>B0Zftjv_-FrW{z=~F z`R(udJs&x86ApTt0rY17Y)?hPRHXyiUy0YC7+8#HAtSpifah zFw7ZU$G#w?P8SsBz-hbI$%4buO0PQQ=HZH}h*T-d1WxeV@o=!)1R-O5GM@0rnQGRC z@MBqFBVyiP6v6+b^hsCr-JtZvQu>bXxrM<%_8?gc4Z~-XVh{vaal*1R8r`m*Rh!57 zCbZil9~)q$AnGf4T9A>gS6m$1$c_cAEc~1k<8W__? z11}QoR)8_Y8^R-$k3s9~3AfmdpC`9jg-j;MQegKXaYfb0EdM6N+Z%ClU3?n*~k`(wKF)92nO!(_l`h< z8Jfts-j!YomwF}z&(SNm|6~iA2M~@s9wJ4kBaWcZkNPi?T%sj@-|IeHIjRkM&jBBx zNx9FV)1mX0+3oB=le$|@>cGhs+No|cbc+~{xOR#Ide;AeQ!NYjw|z-+$K_7ROC&M$ zI#J<|bZPH04$;o4RcF$Q6NHl{|1KXIoFJ5&+hO~vs#&XOql>xjDubKY(s3D&#l21l zv!J};J_Vm3J!3)rzT?VRrk94E^NNa~3=C!# zOGC<`0-DQ0lkKYt`5yJr2!rz+yaX;K*E4c?C}4Tqrqds}PdeAF#KV z!bH%W*Bb~;Cd;JkLxCi|n#^Ka7QuFiEVH$1V%@-FM=$|Dk>oY5ltAVkc2=axd3^| zqSj8r`Jc2Ispm0%gDj4vC5<(3O92176BRQ#ZW^y;RRoH4^^RnV@TYMqzq=kufTy0w z{RL~lLdzd|9uTLA9E8HMYH4u^6VU_n5fjn<^Kz!HR@{RG^hbWy5C-k1;};GF!Y=3T zcWIulrNWwz|6zUnJ3F-e=J@cof~`e)eENSE81*8Pp#lcMcshpCEVkRYnjlQ)BES^dg< z#`Ng=s@>hH@;UXPKmlF|qYn!^Z%_p&jUsOF6CG!g4oU?sBE8u{4&kRf#Qz}#_t}rK+s`bu3m%NT2fjf6O$b{NYDnDboD4%| zj6+bB@$K%6n#ZyDj@kfYeU6`*YJG`58_ak%mReUBuy^J>i%Dht*gZd%G5CIoYvw8bzPYx_f^ij zud}n+3zvI{Vx&kWE=sT+>e7*1wdR8)B8eZ@F`b+buh#8vcAd&b=vewHvA)(85{w5Yh*#^@SyCQXz$g(K z#&w3df~!$P608eEt2}5=;ourVTuR#0J{=0E7_KD6sT|YLDSK&$Se1aR*ez5C73L0m zDrZmiS#ep9UG=z`BklktMYfR{ERNB|R!4LK}a_nwH zSd&uHq_V&L27g)7B((bsXPZ>LOmK{{k(5b;U z(4LN79p_sb%FTHKCw5DsEUVXG&%8!&JzU#+_>aAZJ$8}MLUR3@4a1VvcJr!V4Cxt1 z4=!jb7phCPI|Fuk^}@K`dbD=2x^l6tr0xwqzj|>_&p1}L)KYotn5o@$I%W$qR!k(qC9JHz4lPCzpE+?M{;8L(n>RLEsMWqZh?8cCOP)2;3-AFWjgXv3j6#gFVKfp6#7IkGZUFxS9Lw{(AO( z-rxS7=lMO)zVFT^KTnRxHe_bCnI7fV_l3q@Yp9odykCVt0D$KgBJLq@{@tEu$cZ)5 zvs_s4qdfM!+}SS&;Z;faCgsQ{tT|*$@xHsmjxT@zTJuMV-RlQ_(VuPDw!KaS00B^N zh6Mn<1_3@nfKn6iXJ)zY$uR&tXIb!1WR_2+10j?PT*e)+F|LGGEVKH<|H+iot~vtBfikJ0$2&W>}(U}PVgu-kShOH zAldGZ&B=jmcgB@%tR{er!)||b(Gs_211sfg85@e2Qf}I=*)1FAweI)X2-Le}ky62U z9?0SHPC4BP85=c_02>YPu&jyEY&M_>a?l!oY0L#dTm;{z+Ufz)jbXU-#~dSsy&G{> zAx*-8fQ8h&s+kkFuDwUT;x-e#P$NZIO_zB&P!1*lERTCfo6%_{B(zNJMv{?k#v*wv zTe3|b1$_#Vy&gRDgzj;|vZ%YJ*0oFSSKUN!#&K%Wsw2hNzX9LFYi(%{v9O44&><4U zL6$Q3y#V@kUgR)nZ>N|*Ga(1fxUGz63$PA@*Iap?s85dqW)nVk$bg#K>Ag7Z6*OK> zmQeu?!C(~6YT?-u1|_teU@V&Wa;;{pjj*2FzV+vaT=>P z$-xHAB-KK`leiU%W^)4lCniwC(=1u|02eY@c7(UHns`sR(j*n#^=lg@Gh?U`X(5MP zb?V}5%xsCc%x!LmU{WSE9O(~VCv;8JWdfsf(UiR;@*aW&-q2c8IHB>lMzEV=ZCf~# z{jIty69udTNoNFriE3w8xUJjNf5txVWN=LMB~-7C{y2jWOM&& zPnDTg>(!=}^x1V!vt^Ni^fkOHHJI_(v*#A&lRxcW6i!}O8XS7zdqK$W z`Xl!UT{L7(awHSB_k&BV<*eV*VRPe-P!~s7e^#|oWriWijsIHn|2oDDjm3}9WM4nt zTd8f#2+jR)XH%jYM0}h) z=}oRS(T(nG>dp~FdSl4Mqn5!=&tFaURgZ$TUd=tcs5awKfnte+ASz{;NE>}<+3JL! z_=Hw=;J9)>MpE;l^Sb1ze~M1k4lF=vPWbOWD@qjP@uf+1bWiuC`;lbod%?^=>Wnw} zQ!$wcrxN2%dP-c&CjL#%y{yAs|2WL|Z`Og)j_*i@co|7EM7Q(hY4>%}X*CauAb;i! zgqi$R!e^#VBJJq2zOHA^&K@#xXoXSO&-Zh_$r|qZVVHksSRM4^qVrvT&SKU@{=ZyZ zcls0@>&PDQQ^>qa^PzBkuo=I6c(sFdL`lzpq!DK3<}=>U{CTtMXa(>K>r#egy3z6n z?>MRSX(rK{T5U^=%x=-<*|d{`mAiFLN;AE46;U7)GYU-Ui)y#STA>EjX%vNFiDY86 z_N4-53^#7GE4f;=MyxBgCx3nsjAA7ghWS9P?s^O%w&)Bt-OR&r*q@n?G73Df>$ml| zb*$14#ib2B?}rct87t9BD{Bj3gOEV=graC~e9*sGO4tf1!})53zFJ*$9HRYZxo^a<$i3qs$hFsX~+3L1K*b7v6huPyy1X5j_&yIJutS$bgH{BjHqL6S7?7+Nu= zFpjCU$5Vai`X%vifRPy!ydlDo&b=}(m(U9rWne~UzW9ySu@mve|8+2SbX`fQaxw8E#*y^z3? zQcp|ZmSGgkKOCje@&^YMJ;M+V2qs&KDw}!OTPv~-$T?c~D<%tT7f_wpeS{H;C z)ZDMtQ4KT{7B*PSyoc)Cm#6s~5l2qiXDessD<98SnufB+tDJD-G;cB3@vPHL=k2Ph zsOD)&^KH9mA?PoNXue=4s?Idub~Mk~BUkVxr<*^wlU3m@bCQ-v_SaYNJ0n^=c4yVc zEsq>6-> uint(2); uint param_2 = (flags << uint(16)) | 1u; write_mem(param, param_1, param_2); - AnnoColorRef _808 = { ref.offset + 4u }; + AnnoColorRef _805 = { ref.offset + 4u }; Alloc param_3 = a; - AnnoColorRef param_4 = _808; + AnnoColorRef param_4 = _805; AnnoColor param_5 = s; 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_2 = (flags << uint(16)) | 2u; write_mem(param, param_1, param_2); - AnnoLinGradientRef _829 = { ref.offset + 4u }; + AnnoLinGradientRef _826 = { ref.offset + 4u }; Alloc param_3 = a; - AnnoLinGradientRef param_4 = _829; + AnnoLinGradientRef param_4 = _826; AnnoLinGradient param_5 = s; 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_2 = (flags << uint(16)) | 3u; write_mem(param, param_1, param_2); - AnnoImageRef _850 = { ref.offset + 4u }; + AnnoImageRef _847 = { ref.offset + 4u }; Alloc param_3 = a; - AnnoImageRef param_4 = _850; + AnnoImageRef param_4 = _847; AnnoImage param_5 = s; 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_2 = (flags << uint(16)) | 4u; write_mem(param, param_1, param_2); - AnnoBeginClipRef _871 = { ref.offset + 4u }; + AnnoBeginClipRef _868 = { ref.offset + 4u }; Alloc param_3 = a; - AnnoBeginClipRef param_4 = _871; + AnnoBeginClipRef param_4 = _868; AnnoBeginClip param_5 = s; 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_2 = 5u; write_mem(param, param_1, param_2); - AnnoEndClipRef _889 = { ref.offset + 4u }; + AnnoEndClipRef _886 = { ref.offset + 4u }; Alloc param_3 = a; - AnnoEndClipRef param_4 = _889; + AnnoEndClipRef param_4 = _886; AnnoEndClip param_5 = s; 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() { uint ix = gl_GlobalInvocationID.x * 8u; - ElementRef _907 = { ix * 36u }; - ElementRef ref = _907; + ElementRef _904 = { ix * 36u }; + ElementRef ref = _904; ElementRef param = ref; uint tag_word = Element_tag(param).tag; uint param_1 = tag_word; @@ -579,11 +575,11 @@ void comp_main() DrawMonoid row = tag_monoid_identity(); if (gl_WorkGroupID.x > 0u) { - DrawMonoid _1014; - _1014.path_ix = _1008.Load((gl_WorkGroupID.x - 1u) * 8 + 0); - _1014.clip_ix = _1008.Load((gl_WorkGroupID.x - 1u) * 8 + 4); - row.path_ix = _1014.path_ix; - row.clip_ix = _1014.clip_ix; + DrawMonoid _1011; + _1011.path_ix = _1005.Load((gl_WorkGroupID.x - 1u) * 8 + 0); + _1011.clip_ix = _1005.Load((gl_WorkGroupID.x - 1u) * 8 + 4); + row.path_ix = _1011.path_ix; + row.clip_ix = _1011.clip_ix; } if (gl_LocalInvocationID.x > 0u) { @@ -592,9 +588,9 @@ void comp_main() row = combine_tag_monoid(param_10, param_11); } uint out_ix = gl_GlobalInvocationID.x * 8u; - uint out_base = (_1042.Load(44) >> uint(2)) + (out_ix * 2u); - AnnotatedRef _1058 = { _1042.Load(32) + (out_ix * 40u) }; - AnnotatedRef out_ref = _1058; + uint out_base = (_1039.Load(44) >> uint(2)) + (out_ix * 2u); + AnnotatedRef _1055 = { _1039.Load(32) + (out_ix * 40u) }; + AnnotatedRef out_ref = _1055; float4 mat; float2 translate; AnnoColor anno_fill; @@ -621,7 +617,7 @@ void comp_main() tag_word = Element_tag(param_16).tag; 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_t = float(_201.Load((bbox_offset + 1u) * 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)) { 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))); if (tag_word == 5u) { @@ -653,9 +649,9 @@ void comp_main() anno_fill.bbox = bbox; anno_fill.linewidth = linewidth; anno_fill.rgba_color = fill.rgba_color; - Alloc _1261; - _1261.offset = _1042.Load(32); - param_18.offset = _1261.offset; + Alloc _1258; + _1258.offset = _1039.Load(32); + param_18.offset = _1258.offset; AnnotatedRef param_19 = out_ref; uint param_20 = fill_mode; AnnoColor param_21 = anno_fill; @@ -678,9 +674,9 @@ void comp_main() anno_lin.line_x = line_x; anno_lin.line_y = line_y; anno_lin.line_c = -((p0.x * line_x) + (p0.y * line_y)); - Alloc _1357; - _1357.offset = _1042.Load(32); - param_23.offset = _1357.offset; + Alloc _1354; + _1354.offset = _1039.Load(32); + param_23.offset = _1354.offset; AnnotatedRef param_24 = out_ref; uint param_25 = fill_mode; AnnoLinGradient param_26 = anno_lin; @@ -695,9 +691,9 @@ void comp_main() anno_img.linewidth = linewidth; anno_img.index = fill_img.index; anno_img.offset = fill_img.offset; - Alloc _1385; - _1385.offset = _1042.Load(32); - param_28.offset = _1385.offset; + Alloc _1382; + _1382.offset = _1039.Load(32); + param_28.offset = _1382.offset; AnnotatedRef param_29 = out_ref; uint param_30 = fill_mode; AnnoImage param_31 = anno_img; @@ -714,9 +710,9 @@ void comp_main() Clip begin_clip = Element_BeginClip_read(param_32); anno_begin_clip.bbox = begin_clip.bbox; anno_begin_clip.linewidth = 0.0f; - Alloc _1413; - _1413.offset = _1042.Load(32); - param_33.offset = _1413.offset; + Alloc _1410; + _1410.offset = _1039.Load(32); + param_33.offset = _1410.offset; AnnotatedRef param_34 = out_ref; uint param_35 = 0u; AnnoBeginClip param_36 = anno_begin_clip; @@ -729,9 +725,9 @@ void comp_main() ElementRef param_37 = this_ref; Clip end_clip = Element_EndClip_read(param_37); anno_end_clip.bbox = end_clip.bbox; - Alloc _1438; - _1438.offset = _1042.Load(32); - param_38.offset = _1438.offset; + Alloc _1435; + _1435.offset = _1039.Load(32); + param_38.offset = _1435.offset; AnnotatedRef param_39 = out_ref; AnnoEndClip param_40 = anno_end_clip; Annotated_EndClip_write(param_38, param_39, param_40); diff --git a/piet-gpu/shader/gen/draw_leaf.msl b/piet-gpu/shader/gen/draw_leaf.msl index f713186..e20fcb2 100644 --- a/piet-gpu/shader/gen/draw_leaf.msl +++ b/piet-gpu/shader/gen/draw_leaf.msl @@ -251,12 +251,9 @@ DrawMonoid map_tag(thread const uint& tag_word) return DrawMonoid{ 1u, 0u }; } case 9u: - { - return DrawMonoid{ 1u, 1u }; - } case 10u: { - return DrawMonoid{ 0u, 1u }; + return DrawMonoid{ 1u, 1u }; } 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); } -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]; 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(); if (gl_WorkGroupID.x > 0u) { - uint _1011 = gl_WorkGroupID.x - 1u; - row.path_ix = _1008.parent[_1011].path_ix; - row.clip_ix = _1008.parent[_1011].clip_ix; + uint _1008 = gl_WorkGroupID.x - 1u; + row.path_ix = _1005.parent[_1008].path_ix; + row.clip_ix = _1005.parent[_1008].clip_ix; } 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); } uint out_ix = gl_GlobalInvocationID.x * 8u; - uint out_base = (_1042.conf.drawmonoid_alloc.offset >> uint(2)) + (out_ix * 2u); - AnnotatedRef out_ref = AnnotatedRef{ _1042.conf.anno_alloc.offset + (out_ix * 40u) }; + uint out_base = (_1039.conf.drawmonoid_alloc.offset >> uint(2)) + (out_ix * 2u); + AnnotatedRef out_ref = AnnotatedRef{ _1039.conf.anno_alloc.offset + (out_ix * 40u) }; float4 mat; float2 translate; 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; 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_t = float(v_201.memory[bbox_offset + 1u]) - 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)) { 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(uint4(v_201.memory[t], v_201.memory[t + 1u], v_201.memory[t + 2u], v_201.memory[t + 3u])); 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.linewidth = linewidth; 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; uint param_20 = fill_mode; 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_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; uint param_25 = fill_mode; 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.index = fill_img.index; 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; uint param_30 = fill_mode; 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); anno_begin_clip.bbox = begin_clip.bbox; 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; uint param_35 = 0u; 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; Clip end_clip = Element_EndClip_read(param_37, v_225); 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; AnnoEndClip param_40 = anno_end_clip; Annotated_EndClip_write(param_38, param_39, param_40, v_201); diff --git a/piet-gpu/shader/gen/draw_leaf.spv b/piet-gpu/shader/gen/draw_leaf.spv index 8fade688ab936b660e7d9aebe2ffcf337a1d4484..77ed9cf929bbf122674f3cbc92726f53b8616491 100644 GIT binary patch literal 38560 zcmbuI2b^D3wYCq;Oz6G0kkEUt(rc*Fdl@FlBpFFEA(;T7h7y_-=}i!k4kDrmVpkNw zUa%q^X3~e!qLJ`|kIB*V=2ZUCusd&J55o@9Oha)jZXL)%?|c zT~+H^qM8?_surq7wCXYY9kSm>GX}ombm>;FQ zNLid1+Lq8TuRXO(2Qi@4mk~PYeawE{V-7fA%zk@xkDE5Fd&a=5$-Ui^`=<7GkMEn_ zJ8q!wlp)AOFs%oHr z=D11SQ+lVYGh<--%y9!7wi<0XbJIZr+tpu#kR8EQN@pf2Wr1^lNI(6OutZndHAVy0H1-m`D=RuVIJ5+iLrc zOw;?&Ja3MFJIBsyEAVcU`=*(meagLt;&)ct!u3(tNA5LtZ2zg&-dEdG>octTI&iyZ z{W>acue$H9YDaLrT@US<*j@`YuK@)6k~`iBlY1u4pzh+N+J!oMAHQs>#?Ei|w==%& zu`XKeQM4P}fx{AhM~GI=yAQ@#tlGN~qu2-kF_L57YcS5D)&7k*?_9fg&t<NJvBa#dTa}KwdyH7)9hm%Pi@a({p|CU{^^X$_9H66E>oxW z=Z;c$kG34ZL-$Nq<;_0mTl=*R0C4Ub&rUVxWio?2rEmPeq{7!Tsl4~~xq;J8hWbbF<0bXPS{}k4Nso1#fB4{%neWO+ssEHqwA>ps55Wx1Gn#|j_O2k z`<@(8`7&Uf<};-Al7Ow88xN8}U$Ar7I$Xfz+^D^CZbr0xotxp$o8_yM(FWuG7juzZ zn{(09hvs5%9vo}tf!B$_d6)@q&%=C`2ckI?=YmITYmBYK68q3xjHpg~r+>>=-K(4r zwRg@(S1U$yujuPuabI=Qv2(1}8Ew{(<~v=>j~6GO@U!8KG5i;^Ief&;IW4S1GwWV! z{jW8}t-W?8eP@o(0%xzC4bB{&Tjw`;=Cx+HJsxddMc|D8VsP%YOY2crxTuJnd=@bu~ESx~|S)pl6_W zd}~zgbyD{l8rO=|EpYQ~`F6MN*#9EuL~dowQC}R4?{d?-9yiiY^1cb&&bza^89ek( z(oua7-X6!Q)g9oW+=`w5f6K54^zOHi`QAp~$@jzHZNG{9 z?@yBdFqf?mMUM5u#Qg52ujF?RIQiWRF7vw&p8OsL59K#_=ltJhb1=-%4QPH3(pU2P z7&!So1TORYIJ}+T3e{uaq5SII|Nom^u@&F#{&|ExlGmre?YugwM+fuhtUd>x)-%0l z3bjAyeWAghgzw)wrGI+s>>p9dz4%GEJNK6FBQyXJ$OKIbmwbl9`sT2FJ z;Gx)^)oX3u`fCXAN-rS*Pt+kD;^Lu+8g! zI;zc!cDont_G+}-e^|TD>fkmX%P|S1|JHnb37k3j#xTFm>RS!|?FN6Q!GGM~KWXsi8+>kqztG^X zH2Cit{Pzw1w+8=5gTLC~Z#4Lu4gOYxFU}27uEUZIzSJ<@Q7r?{+>RV(>#RmK_~-^- zr@_~2@bw#fvj*Ry!MAMi?Hhc@2H&Z{_iXUJ8hq~t-@m~RXz&Av@s8?Hc;@Q3VYbey zyTMOr@QDrH*WmAK@RJ&RMuQJD_~{LPW`m#A;By-M!Un&n!7pv_%NzWP2EU=fZ))(H z8~lz2zpKITZt#a1{NV=w#4z4bJp#|#eSVm&v-)C#f2qO0-r!F+_%|B-`wjlX27ji( zpKb8x8vOYN|4oCx+~BVa;~mxS;2HbhhuJ!-R~!7b27jx;c_a=yU*>7>`5Sz}24ASb zmu&E*8+@4tU!}oUZSYYIzE*>;)8Oki_+|~hMT2kI;M+C$4h_CzgO6$Oy&HU=20yUD z4{q>78vNJ>Kfb}c8@#u{CpP$`2A|sC(;EDw20x|2XEpd~!+1w^CcM2@=WqQkXX51U zF_Zhp_DtS;>dF1%dItLXr|!Lb{a!E!-}ZZFS9LLX`1i5S>XJt6%i%o}C$`?ft{7(P ztgeCcn;L%W(SGN+9`1K0-`_se@cnRu-_hWAH~4)G{=hKaRecIRV^a5wanpMS#__?; zZ;}2Ei?Q!#N{ykn-ibbkJ^6hb+&?g>*Y{%c{Smx;2l&Y_&O0r9djBc-j%fX!jDC*j zpMK()>HRaOIk|psduLMt+IMR;FY-dS zkB;gOwPw!#ugrPg1|I?MpE=;C+E#uG)_iQw3}4mJI;(}yS|1tV_9^$u&l379jJ8|< z)CqkP^%>FP`b_QaZU2^w7v1QNYihTnozaUs`g|>aj>M|tDV__j?x~l_L-KKLUN)z_`gVr2qNU+gZgl*k*ogJ$#3VA$=@> zJ}|v!>WtbDeL-};yA1iy-x}jW=;QezrT!M4Ul-C3ze!^t3u9+|MjX@@X=!FQG*^qF zwfoH~XlwVe7`Dmz&OG=N-%#7du^Y|%I;}GYOVl=2X3*xoU1V^5hxVSD*ErVQ+P$|} zt+^+b1G~5KGwdjM``+rP)&b||*Y&{V9^GIV@2Yl!4@~Nt(F)Y+r?c95m``W5YlH7G zjCWOsz+3B457_*Ui^kz3A0OVjk2bI>O6 zJ3D^)F~0q?Z&!6CT6^uh>)P{nom!1=1yAW2pwF)Ak@!yLV_D1IS^XT%nC-Tm)wA%S z`Qztpcx$iLPUim7FrUuq4-NjO2LH=2-ckJn&e&S>@aiyISM@r))t?Ua80O*T^Y^Z; zJt5|&Y4IxH@u$ve#TgBr!8zPpk85MN^XGSLbKWAH?-=sE;6vl=toCp40~-9GVZ5t4 z4BlFcz9}9%t(b>5e2#**#!~oj34L)TU@2XCLkM;L<-Tqdr-S4Txd`48K z!(06pj*fZuFu%@fPJ>_A;1>@|2V-xS(CxBI?on9qpnT6nAP!qL8O9Ol-^TiFi$4zz3VyVL4$uV|kCYNelNznbUl81~7N)SmgqeTsTDidXpPHf{0e zz(f6Z4t{T19^F3k{V91EUoG+Hg3I`F}H^y+XgHJ>qFO|8r= zevV(ic7E-)##Zw=6+c6&`CN&{)@W%<7n)<$FXJ;8pSfzzZQ9zO+NeVFIjgotq4|te zTc^-m2el0f&0N(ujt2YinX0x$q4`Wz+osTbma1)EXpU2D=R)(jskVDdTnZ zuV$=OYujjY8wHl@qiqf9v|AHwJGEu8twlYGvK;N!rk0!QNOEvr*iPMI>~*QtjbrQ$ zsS|r6u(8#g&yA@^QRbl--~FT3j^Ek(yB*uA8`}@Pel#|YvA3d5?5)AZRx_4+Y%L1Q z-16O)TFp4`p_bcbBz#9|<5;xsM6KrB%6Fr7zQV_V9b@>uVDkz;0Bn4@@eijSir-Nk z0atUM98GQBa@X~k8mn9E|Abo0I**4N&-T64wpSa8I)OT4*T=cjZz8pR+TC-LsI%vc zq1`$41tYQh!M0Izy{A!^v5k9D?W4`O{w$Zcw%2alQ>hbo7TCCIBiZ$*QI~Pof~VNW z>2Uq^vyU^V(}%IN+sC=o>Ek@GZPZ+Y^QqH^G3V5Fbz@#YotU=OZp=%m6Z0~#G1atP zPHjx%T|uo)&3GT6PCVObx6O6biFZBNHfp0mH&CaIHrLzv)Apfvu{myScfE^kVr~2A zyVx9?IexNXGsktw(evqX$|~TmQJb&p|74Bj>W=&C)aI;hyP7{$YwN=|20vYM{l8UX zW5oU)u$Nr@2b47^_WvBUnsfI&wcI$rq_$r*`<2h7s3p#eU@!e&qMnc9eH{Di;CTxE zHn?oxg|VZsF9h#|8*eeN@jM4Ij-}A}U(3EiZC|_Qo7LQNBK}*WEneGqs_mQ9{Med% zj@aMCn!DeU{|vOZNqPg;#`EKLxMQ|NdmPPimfB})P2Jc}!X3+6l*E1=t*hW2Oia0E zU2x~d_9MWycdgU@)>_M0Z-YB`vEK)uzvi}k0B(Pt-DeTlz7qdz0O_=J*IGWi*6jbB z8rx=N_*HP{!#=cM4L3ja^#2L4i>?1BYkxKUAE7o^b^DY1txGNa`;KZ(`hTXj*=OQE zf#&@i{x$ggG)Q~jX>IS>WS3{nw%%3+EaWeqoyr#=9Qwe7k=h1RGNq`R8k%d3qTxcI=6_nwzhhy_4Lwm%;6S z1DN~K_e1xG_l#vLcB?V$!wXaOPZYyA2j=28VfUTfcLBfqhHJmL;My-KxX*sS`h1Pa>Moadv3}7o?CLi z=azhKgFBg}-EXPk_UE_MlKU++{Cezu4-40SRKd;9Z(^~_{VrDW4GQl3_+2b^`OXD5 zzTd=3yWhmZwfjvh+%wkiVI}u_SjkUpaKDGe?)mTcuyEt~JuKY*{1#Snzk`+BZ(!lZ z^E+6$&nv%$g}+2Rzk`K4e!qi-`(EOAuyF5Xzk!9j9)Bpf`S}ejcH8?6EZpa*-@wA{ z&+lO2zU%lMEZltk23GRL8r*MTu{&RW4=cIf!%FV=uyDuY_pors0Z!PlWp|EZqM54i;{Ieg_M;Kfi;8PlWp&EZp}6zk`L_pWnd3=fM32 z7H)sD3vPdY`-*wR+Iek7{&1e0NoWh>DYVp4S+|p|Q3*lpHN{fbW3ym+K?;S&zM10nI{F7| zyLIxs1#FCm=*#Eetzb3h`Cr74-$*f*vBjCQ+rXJKpF#4(y&Y^^bM$#6-W^zj*`o;kl0?3}A-&hG-Nd)dbMRZE}ufHUVl*Kn(M$WBu;(Z)!yxFXa z`~P9Enm>EkUYnZjv+k}@?y1MXk5XQwZDM{JuAUqp2Pa40Z9hx#{Z~Kxkn1D<9QBLD zvcJ!R&DYv}_C@M1QM|N&xwc!U@2`N3k$v`6u-X$8_nG_)6k{1%{7ZCyRm>|98>U z)BpFtW&huY+rRp`bxuD3+fJLYU4tJ|Tmx-cr)O&2b@(<|9@~$>YTAB8EsyP|VAogM zPpIYD(?0{dr){H6AAQwb!;IVde_Ds1V{6XCvuNtc;TK@zW_-_q9iM*oDc8q7-9yf` zwe8*a&x4)E+y}n|tK~k>rdICDFRC^6&bDyx@5=G5%h!gvd;AZSUp{Zv-{28p~KG}uQ{{{S4${#46{o2jz z73#lHwB>C4d##y^efnJg2byEdb3?9Q+Pn%j4oAvho7dpB$^6RobAB@}>x}E4;N~89 z9ZfxR_Ajv7i}d51{TuA%oN0T5qUM~56Z=hYGxl3(>Y1~*!D`M~_S`(YKRO22m@X>0 zF+T(C0K5Ne7ynK)b^ZNWM4tPk3+#Pj9BtV*>dr~>wocyjVVezqvCheTJU^Ox-e(p7 ztL1&hpLx`Fspe~aJ6i(m&(^M;KT}(0Pc8}9{$<8w?4`hJet(P4(r{y@zh%H`o~5~$ zmWA6^o8$85EVbmY0ysJNb9|Y@ig4}eZzXWq-zso5f8I}jBf+-Q<~!MSh zna_LRw$qmN-wv$aKHu6oZVxv{&ztnO16(~h?g&nf#xO@U{cN)=x$gp=!#Zf&87xn{ z-N5FnZC7f!>tVay!M=;yMw>o6QL88Ro?ta?dr-^so;n8Xd#ZZQ)xE&#UbeZIywsA< z-e9%+SSR1P_kpYV-8nw{HhlJjtNFe5uDad+aNBB2yZ3_2*ayJX{H~jJ2f}Tu&GYzR zYA?@YZ3j`N^D~)NgE}7p9oe<`x)S}?*X`)ze`EqGvVo5d)l51)|S3c z0js6&Q^9KKdpcOn_O8ikVAmo1bg*+1eg@cl!_Nd8<3OIR{$AxQurX|>uXfjV7PU6p z`COe1F2{HdT+RD5V>}n0F=|iS^T689sMI~eR~Vow%W|$I%>7_^C7TW#&H`sZYwetRb0G?Rd)AplaZQf50 zQh$u%xXssow7c)U*B%0EcTR7mmM6}~!R20j7_R1iao>Fs?B%}G_6drb`%auVp8}U- zc@(Zz-lvbjGZyV>`)RPY@W;Vrtk1yJ%02v9cw%W!+s}cu9mswBdFn4vjIZu|tj$=y zSnJNmBVf7juo;)NJolHi@jdsx1TOFIFT>Tmzn!lqsJ)yoZC|0NIbUL*!+szCD%iE~ z97*i2!PVpQWbNaAh|ky2)aMY#e4YZ^PCf0O1{>S`op#@VtHJ&{;r8OHP;|>XPx#x z1naXfd*A+_0o(sN6m70W;%Kwo2y8zByVqT(pMkyJ;4NKUuJmL6AFo~2PpED4Q)<^v z-8jCl{sO!c#Xg>;misR1IG+RWL9MR+=hSMp^Fc`y+Tciv9k9TCUBx{xf(^ZTl0oT$_LI^jEMs z+wL#a@@@I8Hjj|%Z{S;~?bn#L)yMtgy1IU@r?FhW*T6ebjP)wD-1T$+{1d!4wYv6y zP^;P2_4^mNT)#KquTwH!xjy>4H~$Sjo09c=6RdyMPp(ha<}I*&mh1O6Ts=Od+KQ3& zlk1Zl=Y`)2KY;tg^R5G~X7Tb~ajiXvI??RUe9E(^3tijg1UJSAuyNGW&wSvrpZVcx z7T1MaZLpsO(9%zNE-i?zE&VJ6Hjer^b>9nv?O&UHo6jQPN60%qi^BC&cisGbrCM@b z99-tQ1YFJHRnEzhXy)qtmUFTcy0+xHG}t)m>1P>m+0U|YHA{0&mP1QF<(w>!t}Xqn z05*<##=9ce{tcRj~8sS??T=0_&r0dw<8PwhMOm z&gx*ba*o!3tH)=}ckx*ZO+DxT+F;wMr`%EK5`e^Ejy#d&E>gSNR z_vD6P`_y((t!)HejiSx5`ul6O^tB1NIo3_l)SW-$ZU(kp`raI@kGe7a9k^OzZV7IV zaVs?S#M~NO#@q(3kGe7aJ-S-_w*}iz`gjl6^Txj+(_gMn{I>%eBmUci^}nt5m+KSz z4q*34x#l~<)f00kusLKNb_VOCo;bUJjZ?1guG%Syvm4ml5@&a?KI-Sx>$V5j_}b2{ z*I`evdiojzc1-@=n{y!7C(paR!SD9G+Xr1+o_G6#)x3=D^L9V@UX+~wa{U~G&%5`6 z^Ss+1EZ64q?m%#!cL#vw+D^lEFxZ^U_aLx5&$~mwci>~c#WK4OfrP31DO7c_-ILyMN2s12%?zn1_GYITo%SpK)Mg$8{rF_Cmp{j7n@M>; z#W-T)xVL73v$swK%iUYsVml3d5Vg8vIE7lxd2=tF0anY}yWh@)tHv}G_wyf)UU^TC->-q56l&q^BK@j~jH zkLTleG2Ag|yNFtzb-V=pFh2HU4BP7Cy153&BFYEFlaMtBkuw0w}yugRS=4DQ|QOmdGJ=u4&+rf`f z8{2l;U5Ct{b-6CCPkDcNpC!kSU^9ozk@xx?XzH`8_U{bu1gouFw7UzQcIGJ8&wkdx z&vmrUez+UkD0pJt16RxapiRy8qrtXco%&7Qc(Tp`zp+pIC2DN@B@4V1*jUR@`}ex`^*PFXHP`q`{JW*cqy5KsNf6m8x!Pf&ZczS&TJg`#GTV&_V% zmOj4$w$Jcyf*qIp+ut{QoAMOJXO_O^rlwyx_nx8t?#pX=igV}R_pMN4=VC>Q{rk@q zW`BJbtWUYWz6V!NPTvRHHs4%-2v+m&@v>i@0ozX7fy|Bh{D`9NI*YT;e&1M);x&q5 zt|O^OQ_OF5>ij+SPr>DS{-WS7fnTn9#`80{amw-h9Il@4%Flw;oKxr1KGZVa7r<&6 z+poZCp2b<0UxVjT>_>ar{syegxaKO4?G>U;q7a@{H+EVDl*F z|Icvs%>Q4&YEM%fw>he1p8f__%l!WxtX9tdKj3pI_M<&*Uj=J3u49$Q_D^sb_g`?e z*C|=oH^7OjJ#GID)@EGiOdi`?;4&_^IcqljTmC#~iK{(r=LKssu5&MstrJ|v9RXJ> zzw6HjpG!OEL3`Tz&#h=Pu4^cdZ9#AucVW0%`Hg)M_}n_K_Ox9Ttj)OY1$k^sfSvQ~ zg(czc4|CIBu21}z20NGez1lKxwb+*hJ14O(2Uqi({MzWtgKeX3pDR$WNU_g6x8(XI z=1O4aBQaM7tHr(w*m1=^60DX!R|VTfJ$;S>+h_K(Tt8!aua2(m-pAtH$D3l?g5tFq z#rt?8>dh(M$6He8K3;=%x$iv7*JwdJS?FL|N?gMj|$F>o;jJq*ht-LQcfzPeuYERou!P<=L z9LQta0&HKIgDt_%OZKB&pZISDHpUb6d(}2zHUF7E_ru2Yzb#n(EBMIY16I=~akc{+ zCvCO|tEJ5jU^UyguQnmhj$mVHGhcrm-I-#(*~@bM?8m*cOKo>wiLpc|r z)z;M8QQTKMv~c^MciNqH<-Xdh!S`?QgB$#af*%V$uI8EZJ>d3R?z=tF)U)r#fYq|^ z%tYEY9nU*CYI82+nG4?~_n>&~M#)_44t6f~q;@VwQlCY;axPAT>sQXj>1gV? zAI|`*EkyC%e--@C1gkq{W6J$5U@m8aSF8InS9x;vdAT>mYYZj1?gcj2eW=aV_uUI< zSLS*yT)#5c^U&0j>-k`{E(%t+Xq1y$P;gnd{AH>dEzkV6`&WTj1*EYD{@@{Seq()y-9& zTn|S-lHzp)#r%$<*6;R)-|;$e9Yb;ckETAhz{i1Y-%aiKwBJdLa(s8eZBvf#ZZ!2g zbMFBgtMs`SO+9mPAJ}&4X?H)^{%q$w$n}f;0kC~9MV+(hqhR&;JP5Y$(&uAn>goF- zu)N>bXN~3h$sebd>o5Oo&F$~)YUo?pyu_MJ@tQ<2 z@A;|wYTQfxz8X)YK9OSFDb(%_bNUi7-5YY-exu;pzt!O1F1Y^B6kPuwH~3EquK)7| z*MDw;Gnhzg2Mk{r8n;zV!FsS6*`e zedY4($uHC1H7M_+ub`>tS^NZ8?P>hnx2}m=_S{#&YT5f=1FPlT^-1v8DXy>msM~fe z>ZfX(?@jK-r{QYa&GVb!K5F}Q?d4+E-G0Qbx$W|8`a9sIX`{b>+LP;d!DX)BgRA9z z*j&F)Nv`V2^#`>rx&9EYmRx@XPOk1dx%s*0>__Z=vt4riDcD@~*H3$L{TaB-_2+Q4 zxAFIV%v_(PBvLu-ewd`YP+{`LYJc z*Wdr1M9B}VdG`3SaOc!#O!n$>XzF>EEe}@9v&=bGi~kB>wLHsK1gn+LvX$Vjm;GdK z`z%`-U0e3!DqywjvsJ;azx&Cxl#5+Q`w_c_w#zl$#i@>#Yf zJh|rHG1s-wwI$cJ!D`8MU2t-BAIZ(nJ!C&(_mAz8>jq$R)n7mD$#p|;nd?SywenfE zF+91J&$3O>wI$b0!D`8Mb8vEX56jKZ{cAtsEvapnT(<(7tN!|FPp(^o%UrjCtCi2P zZQ;qae3rcjU0ZV94y=}3cK|0>?;pAOdGFYd*!#wI$#rM2x$3W<_T;(?xXg7|xLWxv z+YN56<+E&eH1*`U2Usn6jsYi6?_0UKdC%I1*!$IX$#WmDdFrp9_T;%QxXg1uxLWxv zdoMhFdhg5a$NOIY+~)^@9Z&fzI~Yw}f6oPZ@;U@;UdA+*_U2hV8@tyTl$_OP)_4~6 zSv5YL`fQ44^*Pizs}Ch+&L+>U;|lKC)!pDH6kPv_1=qi?!QWSK{ZA^m{xceUpy2wS zUU2=-Z1A%RuK%2Z>wjT`UsQ1YFD*>O0UdY&CefYtKsaLv`S$BzW7<=Jr*Sgm|^91VB>G@l*E zpli$hcq~{g_v7(k_oMsUy(JgBuk1(cp0Zu`eGk}uufKlUlj~S;nd>;XTKVi44^OV; zv!fSXTXLNMR!gpvz{%BpEjK^+wEc+P&$dghCxXpYfBm#4*U8{A*C}u{pAq@lVk$hj zs@vA*TtB+D_!&@ro(wK?Jq4~-K08i@ zC)e`XF$-N=ay<>KmR!#OCs*%hx%qi7+mG1$*mlYFY_PfNub=khIvZT(dJbHze0H1* zH`nsnaUPm_@;o1`mOL*2Cr{52xw(0U*oWBj!*71+FtX)Nu@>uPYB*EMjp4vPKxbL6#P z`|F~%oqO^+uzKe2dazpN??!OOasyayZZ}cehxle{+a=Erg3VKZ{j?{~Tfk+WAA+mh zN=cp{2Ak(?)XDR9uzK?R2v{w7-U&{gcYx*Qb{DmMi0`JhUGlsKY@Yh-r#*Sz3oi4# zAFg&EC3$`nY@QELC(j4L>dEtCV727=ad7f{2rM@@p9%IM{sgt{lIJJE=BdAa+MDkT zmtgn0fRcBH3u}BH^+h$FLwzyDcZT;<=bhmZVwUd=pD(!2tS>hBmkRDP?&}5D|LF$* zM#1&}e!=zsVS_(YaQ&Yxxc<*I`11wV|2GBK|K$dMrRKR8KLx)Czw(*;D4KenxsQR> z^2~M4)Uszj4OYuD_i?aV`ON(c{NlQw^5?bBqHD{(`5ag+`{oN^*VFys+Q`MOi~WdQ z6We9qd>QP%(O*C9$@MGXGS?^IYUMNctMKGnK6Af@t}VGf306z4Pl1!G`$BGh?g{%5 zyB}YkLF zpZn2%#LrOMF1h{)Y_9t2r#-p;7+mK16S!LW%>5}mxt7n|pP_3@u0IE>CD&holdJcG z-2A*3>__Zub=khIu~5#`T|_7eCECgH`nr+`zti{ux+XdhzlKiehG--69kfBm#K&$w%_dwqbCGw#Y7Uq*dZjjy1-n&KIEEp^Vg-w|`Z zntMk4z2FPLUv2Q$3a%U~d^5NExFDIR!go6 zfRn3xOm2SeFZ&U<_7*s~E(|tT{q@tHTo(bCxh@JGucyAL#@Q2r{MOnZq0L_E(bTC@;+T2O+EMN3ShPJXMq*r z=APdntOQppe}}L#Ts?c;ag2mJ7tXC?lZzdf{fHftZ8PVi!Opq<`e{%8tAU&Oua2gk zb7Kv#n$HgZCTS|M)`Y8PFRTqtUe2N1T%13B^_xb!^}w!^Hs?~FJl6-GjLr78)85=0 zcVPGW5G8x#wi>%PK3rqh`t}<8E$t%|_r{&n*&CY?zuX&JH29VUcW-Q0@LBL33T_`e z);x2&0o;7by|E#hdiKUfV6}2@Yz#N|?2S#}YUSS86t14L#c^y7zm3|tb!>96 z0)K>Bn{z2oo;!nQW3#>Ov^V$0{n)+kqGWH}U1RsgJvDZ3+*{*`)b~-`8xK%tZ;T;+ zxi|K1@O=vI-Z-$}z3_tzZXbu#JafAX+afBm#4|NX$t{P#yw&)#@1SS@>FJh2Xd zt7mT<1WsPgq1;@YKYjI^K)XZ1u9G(BQl2~y1N-+swzr-3=H7T1yVpl4*&7em*uC+w z8oM_hs_`7^k5k+mpP9LTsXImO)hp^_9J#ow#}TM0CsQaub=kh z-ve&uKNd|rdt)3}Eqmi!VvUEZC$9twik*7a0y@^UWZ=Hfi)tKa3cJ00wNYjaNI$@2{GHP~!#JMCH5 F{|C%wd9MHf literal 38596 zcmbuI2b^D3wYCq;Oz6G$P^3ujRUov01nEtIVUkRefk`GL6MF9*MFgcOy{n)g;zdwV z1Qo>!qF?|!A|g@*<$Iq0|GYDM4%~bFe)nAW-S7LZwbx#|oPExm8K7g%Rp+XzIjVW9 zxvRUns@AngH780{%~y?T)nmu)J8r$1gA>==WYZ0FSg`77`RTI|?QGUbncCAgO~VvM@2UEuvvgduo>sVnC~JB6QOG*m2!s_uO;rxShKvOrPF8bMUzS-tPXsX}#SO z`)2e`80m^5i-?_h;8FXO4I z!GT#5rgTs3ox0}C!5OnA46fU1wC=18H{5#GhFj~`SJ!Ee1cjryt+0YDw_U{k>Crrw#7YJE=VmHLu}u%wH`Bul>6F zrcLZUHe(rXJT#y#=6@OdXAsb_EDLV6q3&X)R-oRcufKoCf&PIRtzO&xs(TIh+f}Uu zAMyo{;{VjVW=tO6(>=i;#+|=f1#QT;dq!{1#3VVS(pxJ{{5z_ZjnT?+u4=Wq?YO>a zV`ubC>?`L(o7Zq&qpCIGMlJkXdqCYQ<8_Csxi_b?!^7R5(>FBlg{rl|jacu=zUXJf z7(Phb_83N0>(+60o7ywE7cXclCv~sk{JSc;ZuQyn1-pOMyoPgcYvw-OcK&J;xcRo) zz9-Z4K0MDG62>iWpN#*ZI3*4q1O8)|(#qCx1-BoQ1uD9!cJ(JsOq2@J+U|({_JE^~C@=WS3PO34~+57lqQ#E#eyT9%6 zZI5-qYUiTe&<-4t@OwhEa^7hu#zNI@jTpr~_#Y!V_FaeKELiQ)i1XgHd;eU<6*2ZM zV(id6xo_Hg7U2D3>@^f)foi`Z#?I3w{;v@n$G$`DyQ%}gJI&}hW{-hs1AP!{{BPH>N)b#cBQ1ZVxar?%R5_f6y`?;AX>b<1{f!XBpoz|`?R8ixI{ z+{3-HzgkA>Z1sbpU{K-*|SaIWPST@|eDfgHsA$&!qC+*O$j^(a$nf zKYXa4?$#426Eyr(%VVy{y_~S2_Uj%!X_Iv)IhpGo-M5jCLPpn7^-*WujsUmsr;h4K zaQmJdRrxYtoaQs6^^$%?+cU4PQm%$6a|6p{(y*I~JTO8TYz4Hm;XIz(pGpyG^oa!%w{#vJvEtKsc&tW@0y9?q@U`Tw^Ji$L#x3z_fd=sWpd2Ts1%gUftxfVcBquDSy}oNxP0 z43Q^=(KTOQ;7Wztlw}O-3ZQwG$+u_OY9`JB}LwC;qZ8nF(4Bvp}_j&qC zeqR75zb}Hz{O*Rg^IN96A3U62z5D-vvn#gZ``thH(ns>T58TeHqx#BF9-Y;L;ORXx zdZtqQbKXM@{wREp-l+pKT4(>LO76u^!ri&Id?!(#qt<(8%<#v#*14hXa|Zub4Cm`H zxR>~8YF_PHeg7KFYvBfMTUYfQVBZP7(|UK9)p`x`yi)U;sZMQ-*)64w|DdhbU#3p% zSHZ)vJF7R^y!GRf1WNhUpwA;1OZjTY|I;(x!d?>#I+q}-NqdFAc&SzBB5B5E@XWGQ>;r-WHos3?t z_1O)6af4sp;Mb1e9o0>69{;WRcm$j|cw&TKXZ527f3m@!YVcn+_%jXue1pH(;J9|f4gRGDzqi5f8^JrO`{7x;Z;Y^YR^MvyM;iQl4gP}$f1<&E*5JQr@TVI5 zxdxxz;LkVs9~=D74gS&y-ch{_&)DA>Ve72kYVdzH_#FJuR-QK<4c^({^ELPa4ZdK5 zFVWyjHTcpEKDxnIY4BAWe60pwyTR9K@J$P164SrOE&us9)20y;RPi*j$M(~d6RCs%@&fWT5&gA~? zvHb(%d-`{qcJ#o6p25C>X}j%IzZaZ~Z~MKIpEbcFzmIiR7d2vk0^T!ua_b%J(h;`K z>Iyi&so}RC?RSnV;eKcG{q6dO?~M(9OM~Cu;CD6n7e?@|>Z|aXQ@UqPn9(yhfe&te zi}ZI`jD4IbHHO}LCwdrr@_Q0IFgT^x_hR#X8eZPB&y3(5)r;^M1IOSys`Yy^`Z;i5 z#*t%Z49uGDY2gU)Aq}?@78Kw8mmwuFKILBdSwf$M&~_Y{ zHmPs2KBHP(pK0B_?cZ|oq8r_DP3v~FGkbAIpR4815m~jqjq;nW?!mtP-d3Bo%^^+c z?VCJ>RAWp$)6u=V{^eVAzjyIbIr32dqtK@hPUw#z{m-S|&MKy%HgjX^;X6bO>0=)B z!5KZ%X4Z!2^P>CRW!Qi2))?nQpU4j>^|$c+x{!YOjT8f!A3NhS;*hpLOEa_Kxmpmd z-EUSwTf2{iu=VFV^UzOx!)+JFZZz-fw9XtXQrlRWA)EVl;i2^%-g|0Z6Igd^_uitl z=AKvHaI`Ot_v>jtMy0lu4)JP;FP|Ztw61QI;$N=_;gk~H~6k2 zIKRJux7MQ`u=yPxjl)SkIlOfrbym~DIRLcJ3LoF*C)T_@_O9x5_|Uw!^6aY4MVrL$ z?D*x!#P-j=UDai1?X~l+YtP$f)M|VkcxuleeRfs%$G4x4Wi5MW^(>k(+ig3m=itNh z*Ht|aZ|&9E$=v@m!l$!(rNLiq@YhE0j_OS~V{6UBTO({;)jRN3e>&7-=;Y?}_pYry zA?BxP@o4bGV~=aaSq(gsbGWx2*G6#X&+pjgym>g^G32|$hsW7j?bYD>H28iacvp20 zytNj6Q$2QCF%NF|90G5RrSO?N!l$$9Z}4diK79n|dpvx+zqjl5w_@#nPZ;4dsyZ3o z>bG!o%x8@7>#WXg@beq|f)Sj(2JhvYLc8a7-=7-cGphPDyw!K%Xx~?l@awFuYw#Ni zJ{JG)FsGiaYzKY^+BNjuX;rvaG|zvv($BMB&2x4v`{Yq-&wS%PM!gEfD}3cPZSiNp z!~JqU)V8Im?K3=i7+)>%XM@Z5a^stm{rR^qp4-dU+;<51ay9pFV&uzG>@UyrRpDOI z<^ZcXZsY!icINGvcZKI!p1i{o$Gp|}U#q>(`7&?0d26@NJf~NMd#U-1@oH*iZt-*c z`nB_Gw>7q!&#CwsQqAW|G`2=dTfER5t9}`ux%kXgb8get{?tYnn$KCa)e6mLtlFA| z<~pdYQ)uR@#&I;%htE{CO$yCts@fKX=Cf37n?iG(YTFl@&rP+RT3To8c?ldEYo3+* zc)6Bp9WC71dTy3$>2uVzbY0vlcQT=lZ5hg9>`d2YWm?-_yaYbZ?`VqGiddZ6w&q`1 zwdY!GG``Mf?28k@xm=n3VIN}s^;ww+j@Le-sT+Su>cn3PY&>vHj5NM`Pm{do$|9-W+UfHDkHQR;RGc zE#EDv)r_+hwcIuw;VqOF`rkb{k zsf}s8OQ^M}8SmrNiDx_Qwz-@-@vZ>dMr~!#r>WCMo9pfTY5VN^*c`XEo8QMaxwhT) zJ~qc@j$dxr%yBJp^n7}NvI6)!)aL8@KU!nCy5s&XwK;3sy5^77+S>5-!H?Hm{~y%Y z7_t8l>?K$K31u~k{Xa{s=G;9;EjP|@sO?wHe&w?%YKij#*h~N4Q_n^5K92n@@Eiqy z7hJaQ!q`#R=Yw~`jkgfkc%B0p$6{#wuVr7Rwy#n1jcV>W5&zB67Ow5v)%Fc)en`zd zN9=EM&E0Rweb7e3_c+|~txidQ@1S)R zd=!&auH)Qr=gaoW5Q*h_VKD1vBH$V0Ce;?Qlp#N8De>MH@r#4r0`;+_qOD+BT-kScu zUfb+5@gGI=o(_KuJ~s{0-uGPFdtUjDYwi~B1K(xUybpXwm7CXJsY$t|8AtAWs+#uK zYMi#dvl??n{w~n>!}O6_`~Otp*R0LU_bp@E?;_OZpk`ij-?`Krzw5gg-1njI<=|O= z;lo=y6%PTdRV8XJCon-9gi67GDvf9?hwQy2L+YMy!eGhFQ06Ym3VzUtKN z{~!Qsa9dvLh%{3aZ({lX{U#iL`@gHf{U#i{{tp*if4}j@ zF83R6$^FJ#@)sN2@2jQV@2lbV=l9i;`+YV13haI(3)kOoWZ~xHce0ZEt*qqh7u@;r zTUqRKzmtU<-|uAM+Wk%zuHEls;hwRFH@M%(O8c}1_ZwO4J`?;#7H&Mhk%ilz-^WVs zx3QA@T`b&q*A?7nnBT`@|2^^iHWu#q{5BTuyNch&!o8RME*5?T{FQ>6pWnq|x4qxR z!hOd2T`b)G{5BTu`;Xto!p+z3VkKX+!ABR|`SKfC{N;WlE4kmu!X1y_$if|u-^jup zkKf3`&ChRS;l2a-jV#>p_4Ww?Dsuh1;Ls zz{2(S8<>0yr>M{LvnW32UqJg1XQ=1$*|mNqKJs(mYCh|?U-Dud>pLqWgE|UwdC_ju$nf{ zbGh++ceo5Zo?2bI=f9e5wSO9{=KOe`eFp3^;D=!S<@$(y7POyLb5>V@eKuLYPONLF zucdgsh5p&vZk;?o2R6nR>C5Nfbzn8;`Jcp)UrjNVvBjCQ8^D<}pF#4(y$NhwbM$#6 z-;$Mdu5Sh}M{OHz`uL1e&z#>1cFxr^=eL2?y=>$Bs-@36z?pNOaq`UhonYr&U4Q5K z^I+%MSUv;gVxNEdo3DOf1RqANk2Za5r=Gs=20Pxwz6V^c;g{g9p?*Kn(M$WBu z;(Zxxyi-{h_y4_MHGdwly*4%5XWd<++*4lzKR|hbwu$*)aP{Q)AUHYtZu=0$_h0?& zL#~hbVd@u%Wq;oQo3FL|>|4~2P`tE%yS7`W@9%((k$v`Eu-c;(_nG{g6k{1%{2O$C zezQCV&)6OZH^=rpH1+iNeX#vKQqSKHz}_e7pRf1&6JXnEdxJ6lh}z41wEd8xW}`w}181>zsZLww*R(y9U3YxCYv?PEXak>+mF49@{U$YTBNr zmdExhu*S=YPk=zsg?WEJ=Dy3OX7Ji{}$Up)b{-%Se|wLJ@|K& z7qF-QKfu-RtoO_x!D=tyvUqw^Te)v0B&3!V4(Z2@% z2jvxtXTNsydWrgVing4MZ`7K(*r(6+H_;qpo*Q!g(&jC&aX3P>J80^evwwlrUZ5Z6>|L;zbEfUz6gB5eoYu96E?myebe=e}P{;i)yz`0N6278|vM_cxdx^t4et&{h>*v^2zQ0L@6 zo)1ku?|bWw7nE`A&8@ zxvT)xhozYo9l(Q?Ehs(!OSGw@yB5fsK*p%?H71d9RUwfMP6T zi_>-;ux;g@8S8?jbV;z`q^ela^DGj zF6*FeN3cBcb^)8Oww(2 z_p;4}-2jtTzk%>kAl_QSJ{srgE#l% zA!zD3&kqHwxks|!4}+V(w(O;DuzhHAjt`(#OFlhdwd|qs;Orr7@tXkFHl4ArPM;IO zYG*LV@=0K|UW)x3POX;qQ^59Ax4`xv;Izh6n;$HLRM z_Ov|?tSx;X4^~UxCxF$`_bjlQ?Ol_Tz^+61$zbOu{1mYHhMx*H#$G&I{awpxU}M-$ zU+u2#iPYL`=X3Q8a5=^^;cDKe8RJ>-j8S{qo(PJzFyzg z7l75yr`$#+=5ire-F@`}#W?4})f49;a2e-fxSDa2*CpWceE2wAtvpvg0k^F-&xcP^ zdwD)+yOg5l`5;d0Pl3z%x*V=%?9A5{@XVL?wEZ+#o8!8STAp*~D)6Ed&z&plI_<88 zYtOlJ4Op#w7rGX1%$z5m1$&;T=f3?M*tXit;WN}~>F0W|TE=k$IOEV3zZ=2Y!fyiG zF8$mLRx8hkTi}VMJ#B9VYYV>(T*kT`u2$afcfb=%d)nR!))syjxQz9AxLSGtegU3X z+SB%nU~S${cT?X(aopx>Kib{*-fLe1Yj;ksqn0Pmm%-&;yce$KesSM@1?=U%({>+4 z&3z|MoUek*u{;1*EAP{P-z<9@K*ci4=}TAusM+W4L!kATbj``d6e?{DYpQED&eOWSuSYR;F~ z=djF4KQHT_)^ zZECJT=FU3pe*xBKe)hioKLxh`H7VL$i^S1pyHVJl2D{f?r(c7;-sUY`U9R*?{vWSh z)icz#`4zQmr*0hISAPTEfnp!eQOkW7b)2)oJ5#G`f0kO!w%&^`fKQ~j4*p)~Mey?! z&sF{9`edxX1v}QP+wZ{o=i7i>AMNuH&ZS`^gxUR0B>uD_4@1NiuD8_n=TJHL} zf8GY~My;;>O=>mUx_r&F?ibkp+B`pNal+ROpB&vO0d zL{pDX2iQKdesX=1V;9)_e^2fY&%04@HH(+`ifipTG#8rvnNN8Z&5f?@ViGjQJYeIf zr=NMjWk2)5)hw}Fm;EdQSF<$dWLdQI zQ_jh9=-SfH@?hhrXS^$b?O$8wWJPc}Co94AQ+H1Movm8@R|czvuL5@7JnNm~Rl)kG z+kP~)+8FHaoz=l=!GP9 z_WEGksh>vP-jf@E?Ni&iwYDMHze~~PSXZM~OJ5sm4QJuCbhGX3TH#D80`G2*`+ zSbzV9OnA1dxP`5+Y2n$=JRe}aGrPjfaThb$F@J%oXvMXusqMZ1Hd=pW533h^UIwVG{x-W>`)kywt`_m#uohfp$Jxjy=P z4|Ide=iT9O{qww&>y!EG0o!N!yc-W!kIw|KG4i~V>!aPj@0b>-+^XfPrpZl%b)v>fvY{rL)-V` zhdkIuV?`bplxK-r5q|N#MPy)g8m})N0O~d+8LgTGrnEb}C#w zKBs|Cu6>Ll*GIe0pwnx6_LtoE|E%ko@b_ESv(U9=UC#!qd1YPCfuBLiy2|x)?5^v1 z;H>MpV7WP-Po4AeO#Ci{I|gkRP|LH97lH4=$9{}qTYX$N*TDI9-m`9(fXj9J1YGUo z6xZ#O)LyOcMBqy)YQ_;~-7W`b-981DXWgy+)H!npf84bMUJvSr@r}S(od>vhQX$f$yg_w(YdL4w*mea$Q`X^8WHZOO7{VGl$HP_xdep>Zh`I<+p;>Rw~-v z22VS4l41e*GIV|M@eq&$$ICwq1z&-_#FKr@eFMKQHIy-`%JEA~m-Cq6JlEWTC;G}W2M>ddug!LHf4AhgzCm$pw$XMg@zlRb(dIq#2(_2@ zjQY1IYUU_*uEc8T^Lt?X4F5jZak;qlUH%Kh~uTs=Ac7;M{obNwk;&A-RXe)$>LcG~u0{O0p>in{A8 z&N};jV-iwtehD{DIi6?W>iMqxE3le# z>U`RVTE_c4SS@3F5v=A}oOSsvcs9j;w5RRwz}k#!uJYLa2v$q3e*&AoG4z+~qrY?X zXNq>`Pn`K%6TCLXYb}cN_W|k;Qk=has55^rgUk7QwZY$L@P9S9|4zBGz5hIk5AQn9GCJVqXF5xME)std>4k0^3GCeU1j(XZEvPKVy2Yu3X!_ zkHxuJ*fyc+Ft-+7j=+2H;=8)NtH*f*|u*5U(dv~w*o z*45F}b6>6jR?9f-Q!V~$fz@(feh{oy-j{2`XV?8`Puq3C+S~``E{|i4QGz-s<;fbNI&>3>VG z`nT|rZv|Gj|ShT!4E9>A>c!6o;lwc zZolQe+XYQM`)({)E&I-#)H0vDfz`6_J`7eX_ucOB*>ykK({>zKoBPAD$z%HnxQx3e zT&>)9d%<^z^$JL&;2Y|I1*SVI*b`ZFX zdoWzBi;}n>1<$VIYERpbfwdXeHIv777}&YVnb!?gi~Vr0V~o8AtQPxtu(`%Q0jw7L zM6i9v-V0XC-*8L<+oy4DtKI!GnR*JvJ(YV_uD`L}=Y6%^dqA9fVB6MQ^Z&TEqj(Q2 zMZG=6dte9ZoUQ$|^B%~aoeI~l+_Tfr)U#&?z-sQ<%>8uu?7E%yjQ=RGHs{{HsQXj325rMA5R3U%}4Rwe+B$c0;@Y_W6J$5U@oVCSE>6lS9x;vdAS?K zYb+(X?g}>74^x|~@4IKyuFUmxxPE1>XP~Jk*E7Lt$#pgS&w{I)t1;!~YA)x1eb-hu zS9x;v4ECF+*Y1?$Iu2~EAE7o^pXnFVuFUm(xPE1>7oe#p*9*aF$#q@)FM_L^t1;!~ zYA%<6eVt$fI$S8ibG;I-UzzJwXzI!JYOq?F z>ost7b2X+sxqcRGuIlD0Pp$``A4KswkYauZQ|otq!|yPixIRX4{y$26NP!Op+rFFH z@oB$_80GkGhTEnb-z{kBIlpcN8>{rW4NX0Ba68y`>S=ce*#2ziJjnHn{Z6obFGiiS z=`OH(d_E7h@6zWBXzJg0Ma?bP4E=K-+S+kBBy|0<;uejjyLU3(tOjrSGmd8zNGHm?3zW4V6vuhn+_ zoD-_%j99|M`OJ|6+sxw&41|RB-)YZt%YrT>sY#uK$}2{#L>D z|98RlXL^R_Lq12r_4nUtp83-5ztg_JO)eSz_ZN5jjH-7+E({7$W1ou(fuWK(CyYBWQ zcFk>g$@Mq2ExFEyt0mVLz{%A; zFE_szsqIJXUbkIx{T{UbQJswdY!fwd*qKZDhh>tDdh z)%#6ue%@>LBlbSCU2=T|Y_9t2r#-p;4P55>DqL+oiuZ-{`ge-CdUhwz*J?d^{sXR- zJl_B(Pw#WNxp|MpKkdo$ZE%_Azu;=`)IQ!H?}F2(XNBB;JSX(e zeLg4L@pum>eg~Sm{+=W9}-tBs<#*8Y2V=B3!5zf;crzCL`u z+Lk$-AFh_46&3)i&1~h~jv2s5eYaw)PiM23TE#Ep80jq6JtZ%cfo-d2S?L%AE zZE%Vrv^wYz44d`7B!z?t0lz_O{QmmC&_iKaK{gWuL7AcKzL3uBBY;I@*ueHMCv! z*$2SxGyV0`o?KT4m$|M1S1X@oYr>Oj?j3Vo3td}s{UBH^xvm3FuI?kb`MHPeN9_Kw zU2!&?=egs_RxhGt$e3tDEPoLiVa{KYV*FX39eqhH_KFjt;Q`g^fL7u!0 z0GpRFjitSLR-cO9>qJV<>XT}G9QDaHK7slaif8p{)H$mUBxcSg&#uD@?pZg!!6y`4 z|04>ne}98dEx7)J1=s)R20y0Y`k!2I{ZDQ1(+aNtxdqq%{06_E;QC)uaQ#2g;FlI$ z|EmhF|1}MMZNc@wrQrJC*5J1nT>mc>T>pC;{JxrJPaOnzjmu}p!D#Avc6=19mS=}+ zu9iLiF|b;m9fyF`%4f%+aQ9F1*>M=Uw%m{1V71(jJz)2v``f)G7rU?QN9>-mUH1J% zu=`$r{j?|7UT~S~B)D4n?3fHsuI0013c9xB+6Puku1A8CtNU7Re(q`e5xbvlmt3cT z%~gN>v?tdAaGC3LxSG$1{A_U)Jh`gd*5}*|bZyCXCRi=G&H^V_?u zE_q%EHc$Qa)1Evp0+)GS3|A|k9Uq6MPtPK`{df-PpZoliV8>HFJ1#?0*WYtWp1eK< zHZNluOMCLV99-sg1zc?u#s2&`^3!1Zo15Bp?#a)9)iZxrg4HsASA#Q_tH5$|yN236 z#Me^WE_r?yY@Yh-r#*Ro4qWDWJzVWNO7grBY@RnzC(oO}>dEtFuv+rG6`VY80n5$p zHfsA2-%f42hmbRGh9HOcZU0k zS-vxTqu@TXzSZE56x?Ur_X@864;uW5g6sdYg6sc_27jvH`af52{bx7$^99%cj|JEN z&kg=k&2uk)6@D&$vzCqu8+di%4hC(;mNgp z=01k5ExA4pR!gqm2Paqeh1~qy6ZRu^KiDq0{t#@g`s=4Xx&8=T=K3UD?Opuy%>6Mu zxvJaN^X(_-+LG%}!D`9%=iub(o|K!P`_X>HPf^=0xjqdxSN-+Vo?L$kE^~bbu2w#C ze+5sjpZ9|Oh`kSNmt0=}o2&l%X-}>%g3DZg3s)i2BkRUqby!if7!XsB^}>Ow74z?iuk$!RLX$)!_duxc+l=47b<6qrp21uK#=m z*METqU$Ef%FHvy)mum2(Yo0UfujJ>L%QNd0H1(WWe*>%K%yJIYvIkxTtL4o4J6P=n z;^y9Z4eoq4e@^%Zy0+|t*THJp2XBI1AJ^NtmW!Qd`w=_Gw#z%ZV?<(c(wcycYztas71C0Dwx)ROC*@Z{<`%gxUCuJeM`lI#57MLvPeRNfgkD|Vs;<<1wwdaEO-Du+HT+n}&24A({`ma^+0r=Vlw~uvd zp8IrJxcQX#>2hf5xlfk|tCc?stN=In{0?D7xLWx;gq7gx+3SvDWw>+U+&VV7*m2p9 z*fH5QbG{nbIoDr5?aBWG;AZ}-qp9cISOcu)v%|kpnntWO;p*879|R{a=TL4g&Y!;e zO{d+uVAn~Tb16@r>w%BPW_#OdZ|;p7v3q@%lD+Y{8oM{HtFdc+eU1H=b_2z|aT9g+ z#-_wC_r~T8zD2>^8`~EAIQaGjw~sM3&)lvLH=lBEY=EYoy|E!!t=ta}tLJQS9Gk&ENA27?Ho4ew*^k&U**0^&CD^^8zkb@2|5o5;{#&D|XK!o+R?FTv zmRKKxt7mU)2Toqjq1;@y*T=%%IG%PpfN!AI=3L5?=Z@f0wNcwndvkBxf!*t7O7_Mr zHFj^@T4VRdZ8e@ueLKayaVK^5#)pYt?u~H`zDL2`8~YZ#7ruYN?c;!&XKr_bn@_nn zc1Bar-q;1KR_=|laC6Vz*cGl;?v35x>N(3D$L{c3sGVEKCKo#{`w=@P+h)!`0(Nic zub=khzbCkv|6XY7*&BO<)v`Ax5^Ep0diKVC;N;~T%FV_3(^tPqv^x;&I%#t*<;n9P zuz&w!d)sMm?u~n}d)-CJ-uQft-5X!1v3ui-H9m{_Zi;*3OVrsL-NbiJ<@)zD`1pdm zH>MQ)4EPZRw~r%hp1D03Za(GS_$Zos_QuD+YUSQI1a9ux8;8Qx%Dr(IT-`Ov8XOLH zE}UD(CKo#{`w=@P+h)!ufZZGV>!&^WPXssf??qG3-k1bd%icJhSd-!E$*T{XyqrV1 zxj29N>USpX`oXS~Hs?~FJg0)s$7XxmX>ab0`>}g{nUcM6Z;jm>_tp3k>aS4T4_~Fu zemIJl<$jpi;DZHsKOA3hf3`WX;P!D+%`>0V;O0>7hXFM8?1$-Kwd@D;nE`hk&Wkz9 z#pb8K*xdZSnfaOpcU|?@PkZt>8r;m|7&P^)>#<<9tm}ouIu5R$bv*%`yqpWUxi}B{ d>US~iP6j*Q+ME-4@;n8688+M7PJ7n%{{g*Kg4_TA diff --git a/piet-gpu/shader/gen/draw_reduce.dxil b/piet-gpu/shader/gen/draw_reduce.dxil index f1e48e1659e60f7d22e6af4278883153380e2f63..c001e896f2b6b3aad03bd517ce94c82d200bb760 100644 GIT binary patch delta 488 zcmew&_e4(8CBn&hO5K@?n1UK_nZjvBcS}XzF)}bPq->NdXJ%YFc_MQkW5Q+ymJDV~ zhgYZ1w|x}sxUtI5S777R0G9|)pCjxH3<(UZ3?G4dKzs%uEdbQM1}N_2eBhoU4+F!C z%?H`nv$9tMg@sxCCzta!)iaqCFdHX0ik}hU+01g-f%^ehgEqrDh8<1^xDLU22b>;o z9nxVm3}6glH{=Onxa!EmEb35t<4{7QN;8*;grNh2088^r-aYCZ%@C9?#N zIg7N!86yRk4Tl7{$`aNZv>cL+In36Qz;j>&qo4@0f%gQU^*obwnME0zQ%*auI`DgJ z>yeOP;AmoG;>viyc9i{K=z)U;&IRl#2B#P39OQD=`q8-I&;+h}g9!|5oC#AFIB;^E z5z+zbUzl*xfi>4bBf=;tfkA!s8fe8i-huMw@^E4PPXJ~%d=D>d-VOpbS z(?d2B2Udd_sTp$?Tv;HJcVq#hg8P SKbq}-Fk8$}Eda#?0|Nk*ou5(w delta 504 zcmaDN_eD<9CBn%$`qzpr9G9o1Mc;hwYIHTIjFEwXp=6_EIWyza$rG9T7}GW@uw*b> zPJVX(T=~(}KAux%Dx_=s@wr5J`dndWU`SwKW%vrz1L89PX#t@2JwS0M=L7c?c^DYp zY(B`oo|U~GC@jnpIJum+sh-)TfWdNz-E>sI*f(^%sZGKa5ZQ%tP|MKbb#xy z122TfaKMRyyGfg8Gs{s2)&Nx&!vMw*PDh>)hR`JL7S>5TS^)`17#JM{c$S%en5V$w zCd{;;$-h0bfuT`IFHz{si@wOFlxYexRXy0U7z!MYJQQGOmf*Eq0HVg=(7M#1)G z8J;4I6AKmE7csGB1ss@S)o`%Dc|+qto#r4GZJvb*?AN%QHSQcV&2iv3#=s*j;AqIO zndP(t>j8~O;{ZkjuWLZJd|8lkwxf};K|w-cK?7S}K#f2Qs|=4hOSuEvW*#5mH3mzX z%60zBJ*2Cy{iDzmXMG$k~wIVAJs;08uUPG*A~ z2b+To62)8w5_ue77p!orRkNuBk`-bL7CJs=%0s#Q`{+p-( diff --git a/piet-gpu/shader/gen/draw_reduce.hlsl b/piet-gpu/shader/gen/draw_reduce.hlsl index 216d923..e56ec3d 100644 --- a/piet-gpu/shader/gen/draw_reduce.hlsl +++ b/piet-gpu/shader/gen/draw_reduce.hlsl @@ -44,15 +44,14 @@ struct Config uint pathseg_offset; }; -static const DrawMonoid _88 = { 1u, 0u }; -static const DrawMonoid _90 = { 1u, 1u }; -static const DrawMonoid _92 = { 0u, 1u }; -static const DrawMonoid _94 = { 0u, 0u }; +static const DrawMonoid _87 = { 1u, 0u }; +static const DrawMonoid _89 = { 1u, 1u }; +static const DrawMonoid _91 = { 0u, 0u }; ByteAddressBuffer _46 : register(t2, space0); -RWByteAddressBuffer _203 : register(u3, space0); -RWByteAddressBuffer _217 : register(u0, space0); -ByteAddressBuffer _223 : register(t1, space0); +RWByteAddressBuffer _200 : register(u3, space0); +RWByteAddressBuffer _214 : register(u0, space0); +ByteAddressBuffer _220 : register(t1, space0); static uint3 gl_WorkGroupID; static uint3 gl_LocalInvocationID; @@ -81,19 +80,16 @@ DrawMonoid map_tag(uint tag_word) case 5u: case 6u: { - return _88; + return _87; } case 9u: - { - return _90; - } case 10u: { - return _92; + return _89; } default: { - return _94; + return _91; } } } @@ -115,8 +111,8 @@ DrawMonoid combine_tag_monoid(DrawMonoid a, DrawMonoid b) void comp_main() { uint ix = gl_GlobalInvocationID.x * 8u; - ElementRef _110 = { ix * 36u }; - ElementRef ref = _110; + ElementRef _107 = { ix * 36u }; + ElementRef ref = _107; ElementRef param = ref; uint tag_word = Element_tag(param).tag; uint param_1 = tag_word; @@ -148,8 +144,8 @@ void comp_main() } if (gl_LocalInvocationID.x == 0u) { - _203.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 + 0, agg.path_ix); + _200.Store(gl_WorkGroupID.x * 8 + 4, agg.clip_ix); } } diff --git a/piet-gpu/shader/gen/draw_reduce.msl b/piet-gpu/shader/gen/draw_reduce.msl index 550cf8c..26a8793 100644 --- a/piet-gpu/shader/gen/draw_reduce.msl +++ b/piet-gpu/shader/gen/draw_reduce.msl @@ -98,12 +98,9 @@ DrawMonoid map_tag(thread const uint& tag_word) return DrawMonoid{ 1u, 0u }; } case 9u: - { - return DrawMonoid{ 1u, 1u }; - } case 10u: { - return DrawMonoid{ 0u, 1u }; + return DrawMonoid{ 1u, 1u }; } default: { @@ -127,7 +124,7 @@ DrawMonoid combine_tag_monoid(thread const DrawMonoid& a, thread const DrawMonoi 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]; 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) { - _203.outbuf[gl_WorkGroupID.x].path_ix = agg.path_ix; - _203.outbuf[gl_WorkGroupID.x].clip_ix = agg.clip_ix; + _200.outbuf[gl_WorkGroupID.x].path_ix = agg.path_ix; + _200.outbuf[gl_WorkGroupID.x].clip_ix = agg.clip_ix; } } diff --git a/piet-gpu/shader/gen/draw_reduce.spv b/piet-gpu/shader/gen/draw_reduce.spv index 02ebc5dd3dde42750026dc89ddb54cf9df6e77cb..73430c6a938744794d28695d8f82174c31aca249 100644 GIT binary patch literal 6872 zcmaKv`FEUU6~|x5B!tq^0xhLb>J(5pbdgG-g`%|3lA__T)*cQ&Oed2`GBTM7lSx|? zwF-iWE4bAa!4<2b*19X!TBwRE;(m^Qg}5B8pU?Zgw|P4rd^xxG-tS$X`z#syPCWml zBsn3OlME(*NwV*>gyeqpR0mda)j-CVi=+KJ$s|hnJwqYNMrL z1bP@c0xckiw$mWerumHP<3H5oAWDhqL&KG!8*UgH-cqSeO;x75cQxykW}{WFj5a#; zTDP&YuHRt6uhFVCXGV(-{dg4n=%+IlmuzX)C+n^5&GoT_ZZZrwX^)Lf*SiVG++^OY zU!~C+t?#~Yy4#tlbyv*>ubNr2X2Z;y4fyvlnyK~$2p8OtK-w~elng8&&yTtdw~x0?MvsG)$dOh61TZi-F1Dt)ozStYZUXFs&*$T zjooQYF>BVEjVUoBm#DWVM;tXZE0bxCC_S5%#i#4fr}DoB^4Qm6cr};Lf_oO6$+_^! z>QuU}Y=-^GQZQ@oYIo>bzk%eU9ILiQE3>odPga88QmeP>SI@AE*&4HWz==+hIG=~}1St+5E# z>OEww!>qA)scG+fu}6RVklWpfx@&MP`{A?eEF}ky)BBS{=(7hQ+Y>Rr;<*z${@j(4 z!(cbebkp7RKB)V(XSyTquycJngOCT(w}f?UZ+Gq(>a=I3+$?*17N6_ulkLu3)*nc9 z!zSyKmAK!5+k<$$(`ma#>l*Jn1fJ%w?sMpV=Kno>Q?uFj9OtFG z>jSB-PpeYTULG#EJ?Li-t%}o`uHz1RQi|KduEwaBwN>dhn)Ng$v-#Ri)Enay-C1Mg zD`C^UQq7tS+&7s^9q_UG1K?BLS~Cn$-*+>2C-FJt#%lq#>wE_PPH|i8v>sL zJ~GnYJ?lT1uF*Ghly@q7W5Sxv)!EF2SiBjj-JizzE@$(aoz4zLZCMT9RsU{e zvn<-PSuVitI=$0jjT!hiDrOKi->yG`XZxEyKK3T&S7YDho|``_@C|V$R=}OL`?HmQ zd+=_ST=cD@#5ijOm8}B=I9^@DLa%9hm`Uu>4`i81|&(+U?%1Fd zxF5sq(>>PbFx)TV>`N^A@=g}}Qn#Z+s6{fd~t8h=O^Zw^XtWWd(p4hOT^lXv7Xz&oU54UEHLk) z*unzlOvK_mSi?Ig=AK5ZYZogQFz=n%1qIAICw5^0bNynTtH@`sV*Zc|%sVEwwt#uZ z#MT!u@0Qq=1XxTleJ*cEY__ zj<^2Gi`Z}HS%&u@v=boxtdllR{h;7Jq z-w(0KyC&E5za*E19P!s8N1ys{>(O6>?zfV)yTV7b{xz?!r2(o({`lHCP2ln?EG)Mx^&mm+n$-VU*7ISY8 z!`1EUDY!Yr>_h$O9E-k=phsWN!%nNR(>99xO^4rJd_ z=j6Ug_ULzZW0%@f;h2zmlmsL$UC0 zAd7|nYmvn+fNmtG>%RjMH{O}68?UVeiD`4*>X%d7_q7e34vF8C_cN8_ur$U!6N`SP zk;TmC{-~SJyV-^Oe@}7k?x&ctd$$ex2WMs{+_NM98-4DE-v#+;e;r)AJnFt4*&I($ z%Xj7tkeGA(Cppx2LFO`6IcD%C!HZ$%dOGsby|Ma}z>V}|!3tH%uAhU^T*^>?mshn%IkJhSS`o>%?tSHE{a zOCWu;>0_L@b9KhfMjpMq6WR6e$@l+V$YR$*#%mKR?wd1?GyXp0hoRybe?Pi-^#1{5 z`!L40{6o+~kbc%s*GG9T++5cCFx-CSo|BKlKL+_}|2SN`Jl69GWOKwh`6RN~Bar7r z{UeaMj8*U{>;dWZ7)eHM8eBpx&U9I|oRqOZ>*i~H6; z&$_<=_p?WBk3(Yis2n|g5xKXgFM)|iPhUnZ_Vfh0J&8w8UqLobTkOkMk=++BL|y_JqfH<4h`X<`ZutxC&9|`i&f(i& z;?cu*kj)$G`!2HU)6bgf`dHI*<*en#55T_%ImejU_j4@nhBmR{S@m4SUi|>Mx5p>J z#2@Fp+ztOB^d#h&_Z<8P?&mqsb`TP?H|5Cv6XYY1anbKjk;P;EKSR#eKaf0w{B!7O zs0`^F@6|7m_4PCVCGcND`j;VRrLJGZ{0iBa6QGFsHL@`=A9eklPpm=yZ=%*i{-!ty zeIDdD2UU^y7bA;Z0L7lZ1lc%k z@%Qwl$m0K@rgHQ;ANh31Zvhm&o`!6HXTamP&_&3uE&6>KvVLVK`nwoe{2*%@fWI6P zTLWzZaUa$~tD!QaJz_6K*5=+=OFe9tA&Whg=UtC1whjuP%aQeo+8dC?u7GYN&b3|% z8LQ1%WO}x53&EB^erH0?ViA0CK|TxFxttAmF50hxigVeBZcG`9-Zmj`gkpAABb%!X zh0is};<1mLk&P3NxGl)mD?`pnJ@d=X&=9({J+F~>E3$axxDMIcMW0upi%0EOBO50k zal^>gHqP4W`dP=>T@Pt@uF7%7w4Dq2odY>n&-+rySuTS|uQ%rS^Z49~?DsEjllU!A q3AzdD&$(wz-F!Dgb8>E8{q5y;xVnDo+jCui_1kjYdPkG&9{(3USp?_+ literal 6908 zcmaKviI7ZZ8k++&Q>?bGbS(QJ(A^Xw=G$db3s@uD5H| zPJJ&)OVsN6)tlAE)Ns+E507FWeRRg+lFg0Uc&*vFx;B!~O@;v{t&x$*S|Ke7gR8D*tOBkA0mEujKL>@S#-hOP0XL zD--FuvKjUzi@~h9uhph&s*&44omZB;9sdaJo@Q|`M597ataUqU>|JU)uluk^e+Q9Uov~UQ?phAPyX!0^hmX_ylE=}z2O--NF~8!u6FdIg zm69jHc20HD-Sj@F`?aPzL+-G9c_f37`_s3CbzIwO?-^{jrY77hdwd3;9kuaR`+)WP zQ{Axf+ITtcci{FQUTe2ouF<;2`woGpIjsA8bU*X|5x%j}XnBtGSNDtD`uqv~%2sov zp5FETRM)3ju4OL|7u+87vxjEc=}gve2c4ec_OP!$>}741JM~5_jmd1jwqv#W=vb#~ zjC>_*x|b_ilY#ptbEyM{}ECgxXV-{qd0zbfzzawb;7owfV3 zm4D-{@=e@|cj(K}{RW8_bFYo}9XlKH3w;@~=S1E2&Ur4*b?>?QS&;d?bAFLOFwdWu zeVOM8xP7|E`oukqIQtTdzK$Rl`%<^BupdRXk24|jKMn8BA-iut$MxkMEcRs|aWCz| z-t0px>Uc+seW=@qpLPELw~s}T`QqN1&ri%V=hut%_M)G2k6t3yUc_cWfjL((&skvJ zMKS;W2Ifq};yhTxJ1FL!hM#K}8z^AjJFyi7%sVG`P62cMVxFtWXRl)ZkPFN^CU#)~ z^NxvKQoy`hVwV*#*D1E4fO(I^HlUruEw z+>2#+>#w|!{dS%Mkl#`)&N9RN)tJpzYye;97xrbG181kMpSGp=xK`t(W$3O`Odsv; zW!TNR5ZxT&<~%Kz#U*oiH?>>Co?iqXfLzzgf{vQ%=68>!C9G*BT+H>kr|Q-|5AGbq z*5|tKhgjrYlk56lkjv)ME?)~BJJ$8T6x}sP{1)WsQ~%fY=+~nAtzzx2@j6H!@xZRn zG4a4|$bE{~jX7rQ>GW5HThslv|6#bjh`U~O-$OC&Be`r$tgnvj-YsA%Hx=;rAv>tW z#JxYYCyAJOJ^(L4;?`03HFZZns8pN3mwDSP!yE~|^1|95cnYcrp^ z@2l7b;w7)q9(-(Ar4zptR{f3QbC)T1BH^^Htr1bGzl{&?SB!AzX}7$klhKI*T_u{!*f zOw}2Rg?|HCEd2d{Be50GE^<2WCM0gWGgmiW+v_1QZO&VL9kqRP{jEM962B_%ho8qO z4ohR)GqLDrFS3~V+#hxGc{le#OCWLW?x&ctd*^xn2WRF0+_NM98-2bJ{wByz`i2GF5u}eceT)-#uFlxm$fK9{AiMsX^ZkD>ve*{Lcx__EeRIZf#y^OBFH}6^ zA3_(8{y&UtAIA8We-ye0($5;|`Y3zn%w@fg!R=S>Ir#+qe#lSzC*j)Vv7S#Mn(iJ(kb!VPwy_c+B+k$i``l zzP^Ag?pyyH>;4kl&mOgX5fZaU<>=`V2Ege^($n3{mg#>{MV5F1CX;)*DqpzgKW$RP{jNe*_fD*x_-_l)*$~k zQI|merkIX?3gkBvT1%}N@RM`-B={`I_0NV|OKcgW&j1wbKO0#*^ySFL#^25tBa5wo zVozU!Y@D|Ed-_sj@&8a$IeMLgJRkC#3q`NESR7umH%zb`}9Zvcw^&PNu1 zlr{CkUk-__fi{A;4;MnKp#eyH#9oA~&AqXfde|;T7JEF;dkM1GS}1(hA?p*h*CUHv z3hg4!wO$4ptIb(tdbV#1z`VnLr$NqQA^h}$d}ITP+&v|kPt=duCam;oqy+lagY zirHO(Y_0()e6B`B8x|kZOGOx`n(EV zJZirh**Nit+m38)JZ3*Oe7UWz#?~5U4xfC9~?#l7!@Ocfg z-@mv`;@3bW=qji$=bkZj^IZ+igmy#b)!$xj%ys?Lugi7))vwQW Encoder { Encoder { @@ -116,6 +137,45 @@ impl Encoder { 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. /// /// This does not include further pipeline processing. Also returns the diff --git a/piet-gpu/src/render_ctx.rs b/piet-gpu/src/render_ctx.rs index 2c5ee73..31fbf9e 100644 --- a/piet-gpu/src/render_ctx.rs +++ b/piet-gpu/src/render_ctx.rs @@ -11,7 +11,7 @@ use piet::{ use piet_gpu_hal::BufWrite; 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::text::Font; @@ -64,8 +64,8 @@ struct State { } struct ClipElement { - /// Index of BeginClip element in element vec, for bbox fixup. - begin_ix: usize, + /// Byte offset of BeginClip element in element vec, for bbox fixup. + save_point: usize, bbox: Option, } @@ -242,21 +242,17 @@ impl RenderContext for PietGpuRenderContext { fn fill_even_odd(&mut self, _shape: impl Shape, _brush: &impl IntoBrush) {} fn clip(&mut self, shape: impl Shape) { - self.set_fill_mode(FillMode::Nonzero); + self.encode_linewidth(-1.0); let path = shape.path_elements(TOLERANCE); self.encode_path(path, true); - let begin_ix = self.elements.len(); - self.elements.push(Element::BeginClip(Clip { - bbox: Default::default(), - })); + let save_point = self.new_encoder.begin_clip(); if self.clip_stack.len() >= MAX_BLEND_STACK { panic!("Maximum clip/blend stack size {} exceeded", MAX_BLEND_STACK); } self.clip_stack.push(ClipElement { bbox: None, - begin_ix, + save_point, }); - self.path_count += 1; if let Some(tos) = self.state_stack.last_mut() { tos.n_clip += 1; } @@ -405,14 +401,7 @@ impl PietGpuRenderContext { let tos = self.clip_stack.pop().unwrap(); let bbox = tos.bbox.unwrap_or_default(); let bbox_f32_4 = rect_to_f32_4(bbox); - self.elements - .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"); - } + self.new_encoder.end_clip(bbox_f32_4, tos.save_point); if let Some(bbox) = tos.bbox { self.union_bbox(bbox); } @@ -468,13 +457,8 @@ impl PietGpuRenderContext { self.new_encoder.fill_color(*rgba_color); } PietGpuBrush::LinGradient(lin) => { - let fill_lin = FillLinGradient { - index: lin.ramp_id, - p0: lin.start, - p1: lin.end, - }; - self.elements.push(Element::FillLinGradient(fill_lin)); - self.path_count += 1; + self.new_encoder + .fill_lin_gradient(lin.ramp_id, lin.start, lin.end); } } } diff --git a/piet-gpu/src/text.rs b/piet-gpu/src/text.rs index d25320b..a47c614 100644 --- a/piet-gpu/src/text.rs +++ b/piet-gpu/src/text.rs @@ -12,8 +12,8 @@ use piet::{ use crate::encoder::GlyphEncoder; use crate::render_ctx::{self, FillMode}; -use crate::PietGpuRenderContext; use crate::stages::Transform; +use crate::PietGpuRenderContext; // 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 @@ -51,7 +51,6 @@ struct Glyph { y: f32, } - struct TextRenderCtx<'a> { scaler: Scaler<'a>, }