From 9df486ef35255375c04c6ad6799cf3c182fe152d Mon Sep 17 00:00:00 2001 From: Raph Levien Date: Fri, 20 May 2022 14:02:24 -0700 Subject: [PATCH] Fix blend memory issues --- piet-gpu/shader/coarse.comp | 5 +- piet-gpu/shader/gen/coarse.dxil | Bin 11632 -> 11972 bytes piet-gpu/shader/gen/coarse.hlsl | 596 ++++++++++++------------ piet-gpu/shader/gen/coarse.msl | 639 +++++++++++++------------- piet-gpu/shader/gen/coarse.spv | Bin 59320 -> 60516 bytes piet-gpu/shader/gen/kernel4.dxil | Bin 14480 -> 14484 bytes piet-gpu/shader/gen/kernel4.hlsl | 2 +- piet-gpu/shader/gen/kernel4.msl | 2 +- piet-gpu/shader/gen/kernel4.spv | Bin 66224 -> 66224 bytes piet-gpu/shader/gen/kernel4_gray.dxil | Bin 14556 -> 14564 bytes piet-gpu/shader/gen/kernel4_gray.hlsl | 2 +- piet-gpu/shader/gen/kernel4_gray.msl | 2 +- piet-gpu/shader/gen/kernel4_gray.spv | Bin 65980 -> 65980 bytes piet-gpu/shader/kernel4.comp | 2 +- 14 files changed, 647 insertions(+), 603 deletions(-) diff --git a/piet-gpu/shader/coarse.comp b/piet-gpu/shader/coarse.comp index be891c1..fc6df21 100644 --- a/piet-gpu/shader/coarse.comp +++ b/piet-gpu/shader/coarse.comp @@ -151,6 +151,7 @@ void main() { uint part_start_ix = 0; uint ready_ix = 0; + Alloc scratch_alloc = slice_mem(cmd_alloc, 0, Alloc_size); cmd_ref.offset += 4; // Accounting for allocation of blend memory uint render_blend_depth = 0; @@ -460,7 +461,9 @@ void main() { if (bin_tile_x + tile_x < conf.width_in_tiles && bin_tile_y + tile_y < conf.height_in_tiles) { Cmd_End_write(cmd_alloc, cmd_ref); if (max_blend_depth > BLEND_STACK_SPLIT) { - // TODO: allocate blend memory and write result + uint scratch_size = max_blend_depth * TILE_WIDTH_PX * TILE_HEIGHT_PX * CLIP_STATE_SIZE * 4; + MallocResult scratch = malloc(scratch_size); + alloc_write(scratch_alloc, scratch_alloc.offset, scratch.alloc); } } } diff --git a/piet-gpu/shader/gen/coarse.dxil b/piet-gpu/shader/gen/coarse.dxil index 3370f591906a162db9acf66e12c14904dd757988..f71cc0441f5852ebbb2a904e3dc8c03be2439447 100644 GIT binary patch delta 7858 zcmY*e3s@7^ww^o*LxAuIgqH>og#_t-R{=#QQ!RYxD$_cqt>XW2CrjS?UT>Rb4wyA9dz+w=2ao`#+9SM4+F1f3`8 zTQPc*=v_cYOWYvhHYolKK_jD9B#^f_=g8*Mpv&Igl?PbV5-4r|jsEfRim%Bm!p*i} z@?LU8fASK0?cKvWZyTlDz^61#Y}%^T8~q?i22r3?YY4)>L%=cv7>@0}YZ9ig>fn zbNJKB{2e^c@qeP}Sp=Gr{T58nmud|uW%4hDZu1&ZSi{RwLu?u~m0K-12RATav-oO{ zn}1U)e2$kT!^NdR#TciI#e!%`!Y-REBY+`PpavZg-Upv);G^b7&s@gG4EVe?K`e0N zQ*%+92%TojnL zk7tAHBF)`MAxcqLEDYC|5{M|9mG@L`Gcp$$@5vKFQn6il7K%^NJ#wOFrBxU2Gy`I4 zNJm+0FjvJZyKXjh&y-0K7BOmm{{>{X6|#G5f_7SA!$@NFJCaFE zvT8~ufatB%ZGsrO6na5ukH??W=I1LyWi0-E#{7Ik z=-aNwbs$i%Ne6$9{ur)HW~IMPeJ$S2aBM9Z^NaOiEnzETRJ+As+k^T%4`wQ8ci&3JD(g74BfDrh6-4PLOP03xiJy*tvo%R;W2vn$q23FS`sCX?4#_BLPP z^fhAwf3Zc@d8B>x09X-&Dm$TJH9~n?#ii-QYz?u|rIv*>LQRy^!|5d?Z3?d}ZUfp2 zyhR5qLFv-8?ef;c!Ly+YXAbtQ=?|HgD3%A0$a}-dIlSAg{(|Ew2A3rkEhNHdvVhn| zzRkB{`%8u1zEnJ}HE$_vXNm(pB&fzp^nrTW2}Q6nz2qiA<=zja@VeqMTYC{;jt0z) z&ZdPT#9ij?Qf#mH4{Nol_S;M4K-ia*EgE7g<)BNk0)}?x3H;ZotyqP$5gEP2qgX*t zptWau8`G(Iu>U+Xctl3d<~>w9H8@DsP6djMsEpFPqb`n5;FWk(%jgXrux}$8bf4G- zVQ|Tdz2yi>l+l;QW31Ino#VuWycWtQ1SOOHOR@)nqGuKAmBH18{$(2CAnGV2rSST` zDI_~X*`1ipP1&_to)X?eN{F{j!{l?v9K=xB7&U}u7W9!4!tK*A#(l(W>A9^xPH6Te-2Cxn*5 z8$n?!xBifgX|2lImExR5{!kocIf1D$l$*-1424yfR)WCioMSfPFPQ3)Z;;RfI|A^0hd@hhJY-!q636z%^WfkSi=M3^zSkTE z4ImE&gXdmA6fvvLh_w0?Ctqm1Q=7iZCeg+8sfN-^Ah zybmt@G(6N&298k^#^d7c<-zFJ_bf?Dtq+A1g)B)g{cgu;PD2C-(mIX_qKSAFVUG33upi)l-cZ?EDbS-)c|U{!WX0fV{&ux#m& zdokt#OlNgaw@l&lOwtKhT>zq4XNHW9m?3w~&nOeL*S6;u6hi_<9(sW|_L*KTf+A$R z+o*jZqtX~5nePOseyaAYi66+9@dfRGtT6so2vvBX0?!s^bPF?Td|kw#IDF8!@L598 zvv0X&D>-G~`1i%m?d$VG1<@!LjVhwYD0?GHr-4|?OZy@BPuvo;NzJPlsa2Cpc?rU=8PaKpoe zhD||+hg`!Zj^UxdVbfg0L)>Q9ecYzuY((F{e{Ev#-6OH|7MuUqB=TRIB-RLtHQXDc zC9%rPHo-bkLNrQ1 zR>4F5(vQUt1?%mHqNT6IT2YuCsfqwm3nA)GdFr%!-IP4(DI;NwC%mqX>`-$qs^@EY z9_L$KK4v(62vJ``)IT8VE5_PMi25ThQI7}4&(vJg^9la+E52KfH{ytr@M%n~uavEa zOWu0dfZeVnM2fN6uxT5qEN1C!Q=IXiXGz24?!|zOdTq1yF(6DBRxcSXvWc2k$^Oty zxWu7oZV6E)l`WPK?(!*J|6sDjg4cPJCnR%|_SF7O*BVU>PELKVE#&mBEoK+u)c#ML zovN0n3hwYJ|M>^A)!xi_a zF=%&4{mb+*dc7EaXbcp1;mzYps~A>47TD`GED$%zg<=8xv{XO>j4F@n6JBy^XV`U! zZIBM*Z_Zo|L)t2g$TH%p&e&gN{7POl*k!co7XMpH67mw|RwOT%M-K>(ZYkWUAMlKX zy?Y>+U62d@k+lbM-37UR05aEV2ampT^)=D{Fg-)c7G>P9#{xxjrU=0!dTI_CB=A0s z6=t5ow%wxc#iMFBWQ<`qrhZ*17G(%*EyDLIMPEBj7(FJ`c25~^fQ)CJn)_463^{U; z8rd$2tcMCK_eGe^;?hbl8ROQPHsbQBD^WgKm zq8jC1P#53GG$i>n^e>XaKEd{f2{SL8_YD@-pk+8O(1Mk)A%5kog6roO=hr~x`LuE) zt^C=AG9c@>$ZnqxEkbJe<#3q9H>{^7UAdH17;m)~C-USOm8?&OvqfW!kM9iT8DdqZ=?RxaPaZLa@-bijuIY@Z|3`ghMXJ>O7{UVn7Ru-b)WjMYjYR>;JfToU-4-_ z+5rW8`!7CWj%s=JRu(0hmOp2d#!>qsUcq@dxqeYOnTHS~VZX>rzL7q)O$Wo3@PuWe zk@b<2HPA8ls2<;_klOqK;Dxv?nqCRUcb4*|NT1M!TQsDpSV7ENehmJR7b}Q;%a0Xq zIiYZVBx@A&>YLfe%JnQ3Ri z#>%a$Hu}fcZ;G#%Ce(X09jug{U^C4x-L97w)rzVR{;DeJsw#ONkgQgCzb<2i;tlA9 z#3r)Hz!@6g=mt1X#9obcOpT*i_Dy2TJduEVNW|@saDO%0>hG%)ECIUcs-rFC;=HZ_ zZnwmKeSHEiQ4N#4&}?`(fYvDPY)v#*wA10(r?!+!p{EfBbf6n(HucphNo{o^K~Op9 ztKa|9atN)h2zn_-6Y+}kIR`U$3BOb(%AB7KfI=dzNEBlWEt&n>cw&z8+kuF81B>2C zJcd+BV!@*KDBVDpE#CLuSQmaanVobj{7^+WmTV2iky*j5ityTGUecZL^6$cLB-{S| z!kzEJ>y!COHp`AXExVQM&~8z5&uN)9S)8|ZMub~r@fYov6NnTVz$MUVsz-?S{=Z9S|5CxAeF z11OT8LsvfDz4o{9q_?4oMwrAr5lwADKgT%zAzex?`K8Mi`VjIspxdV*mrX&Z}4*LS{>=>pmuz| z{@hwIVflMO0hBiSmwcN(xoS*#WkM=O^*;le9{EmcK{&YgUHH z^mhj1>m_{2GJ#~?3dtoPv{bav6_O_B%E2|2AV3ELnPdqsQEX#g)R1(QZeWY#ZH{FA zR!Ms(5Hz%Yy5AhzZ2=Woe&cmnlXXdxb@^JmPY2+5*evS}sN7Z0uHWxdFX7Y|K1Haq z0{;NkY!OJhAcapTs~l7E>r*%(#u+tx>Y; zR=g(#Xf~aLZJJMBs%JOgW$E$>rMWdWBsA;_&}i}2_@8x_2Dq&$AQLXk2PJ@m(1n)w z*KyChspnt7-#3hYZoyvuAz&w5gkvY^f16TqYd9C6792W=)>xodVeW4DCJt?~m0Gm6 zaW7?D(3_y6`g#krD}wtWJOH2wly8@5!T!kYID7%s;cEN%It}sifSEr2oYYMQpxjGJ zZl{uqa6$~+M!l|`c*Rm^ViUiZ3aI9WR5xEe-K^$dI$56vFW#N^PB-s|ClcrrhrgIe zsLPjg&7MBmK9NwPyhcVl4p(>WhYR{IqOM=!UiZ;*_FKjm2VH-B zVf2g6>(`pEbF?8jmhnj4^~np+#n7%EZdZsl!eAK}1$Di+z`oep^=)%kgccnC&H58{ zU2iV%F9vpl5+&3oAfH$S>w>z8TG_?M?jM@F6SR^Z%lHOeH>^c3&im>y_bZ7uLvI<+ z4*JSbtGTH8YO?vO4DDf6s}c}c<2}BvS$5Iv=Fs66&7Z>jg4tT%DTtL6gLBcIF&E%B z3ZP}p_GaXwgmXFTNPGN|i*a;2{5qvacU__#z(Wh3ggO2}I-hoMFk$q`P7qvofunvJ zPFGD=;%yrKJbXE&pkr6T3AQVh6BuqQ7i?E;sNOnE*+e_MLSn;73o?`2-76c9G$_9x z(0yN|`+icp%V9_fXC#^!8o^M~qTx$g_T|&VhjqhmwE3|%xz>9CiH+_Slx-K3?E*yi z;yb?U+t5}@r!DwK;_X%(-L5#g3$UB(BR=bB_1?oXht)xi=;ijZDR94*30Y-zV$uCJ zV6C5Eu9g@X;9mJ8-xxGK5cF^$sGT!*xn21$r2{hCpWH9F(!&|m&5Q;;dK}aNgysMW zfAu)%OStHl=%VH;MFqNiEU4O#cF-J^KQv#tIXM;d!;7H*c@cCOaJ-JX+gG6Nd{IL=DQ@z*JZ1=PHDxymv?kyEi{2I)D9Ah26nxfA7 zvnng#MAnt{2g(E;@KHdvri(IJ!v>wBv}FSKhTs-_2rR4F2m_DkNs$(h9B?~)&*j+j z;UnF{hj5P|;QB8#*C%BIZwc-alqH3-Gbl&Q9uY-$T0G)>x_h`#JACXxb@}lJfr?1r zkp_=;(j2R2fF|w#%^1`fSyEZH*&`cIKN;RZiAc$fXvmHz$*kgrQJJIFW{+ISjdaN* z91DxOnH_~RWDd213F~o}bdxf5q|B?pBP*_Bk;cMSy9Uor+#_A8e4Q%)YHy$@s!IX- zj0NvgsVk;g(y8nW>Sy4H6J;O;&mIv*&XlJYS?A2(mJ^qoGk-I)_*4ob*W8_LJX`gk zh?0AGK`!`<+w52DR>ne1?cTO+xi>ShQ%`Q4I{V<%P-;LHwT(P&erx3H+4OyGc|T?6 zmE`8-ZpH;-0_}ySUxeK{^U&On6_w=v#%z25`Y~rku2^x}(vN+7KlgWLXL|Y7Q=eUR zL#mGFiV2@tWXP-BmrO*TPS_@=Zu@in@4Iq&iBr^GPkl?5!fTl1H6nU&y9Nscw<}9W zj&!5gZc^jJMDgoOCRiI8|Mg~cyN3Ge0RE%QU248;^UVJn>bDH(EhzLhHJ=Sovp*>$ zv|M#NgV-*u`!t~(?04d=?^#mzM4&2$Z0Rz;xVg={VcREoWc1qpo2dq#j+p`hf#T=H zjFxw9hn}Z5f&!qIlRl>8fMY(G$%$6`5U&ZaQ_G8E#fgmq%6UudJ~*xYGEJV}z9;{} zp8T*SBcuA$tbQb9Mifte+4Z1&w(sjGIQLW5<=g`8iF|Wd$9Lb{wx}T=FDZ)81(c?S z^(=F{IAzavFt>|-h1!(tx)TMh)Pk>;BtniIn{`xX*J(=%7G5gJ###c}mK3zPbGnXa?EsDJz28cDGn>wkk*zK+b91%0Gb8FVmYDYU~8YB!&!!^a=UY z4b*)&7;?nHDLw$KO2om&LpWIGfP({ZaO(;jd=&?yt~mIu37m(6#pO8o3JxYjBRIMx z9!Gmw1GL^3NACx>xk0K^CUgak-t7p`-fMC2%qjpbNyfoy4i3)r!zZo7!NXNJcnk-R zZ8T|4#5M0Vfft$PfrCi~9Lzoiz#;SiCxG6v6-TR+aP+@$G#Z1W#~0w}cjO%4slWgJ E7kU&-zW@LL delta 7515 zcmZu$dsI`$)<1b2l0yO^5(okkK$JHTQ4mBDf=IzaP>_d;N8Tc}(NeUi&5KaLpeSI$ zY8wP3JVkj3NN*DY77=ez!J-$tg0;QqwY_MmAE?!DPDr}ex7NoWoO5RXX3w6Py?=Yp zWN7zX^7636P`)n6CwwAoSQancj?c=8Xu&}cB=own+?41NS+}zEGP`u4cAOH~#v+V{ zh;QH>_Y=iMHV{4wia3u`qA*kuN!)9R=By$^J&ul*au!?!B_10do7Po)M`YpP&MD#% zVqo=u&X0Jk^SF}!8i8zXKcj^uez29i5`s{O1mP(VgndI`WI7mp&KZJ2`Rbn>sSvbC zuJFcP>f|6Q_=A-u5VVZ=o^DEnpw$3J9D^!TxsI@9n2Tk&#F8C$5U%9WD#IjjC7%ad z9^uzkZgYzZ!!fttta{JYSnl946JNlUTVg95Ej&fz+wsm&`Xp4vx0gVWmC|JC$wt6a zX!r(CHj;i4Qdo%D^yfa7xgc!u#QXAJM zB~r^tMKN)NSE0mR!D2yVDK6DCRSjSw7^o3X;ctL3M33xVKQmdHB z+J?g$gb(m;uAS(T(2h&GB%T}_jOUvs+_uf$&?9D21<~xytep~IonBW7!85UdZHnoV z+3zCZH4$FrZKB3ObGZ5l?D^90oW~W0Vc)L|&wt?F>g9)^ z`TgbQE++4w{j9{t4${2@C#VGH8aQQ=sMRqXEvUA1Nd%SqK;y1n4p*~E&?zxFA{9c+ z6LI%BBViv9wQv~9AP5TB<{d>6(W+=C)_OL+k|lmDOAxQ4u0PGtbl@yhLr|rI8sEBC zT~b9Zq7{o>U?tbl3Z==SyTUlRT%z|HsFOIzB?(QL_F#tGkK=Obz}jxX6&{$@MCdTA z5%M`=@b8XEx*srTxjbT6f1(^C9S{M&N#q>27RB zd!bB9sRU8S!}4ku4(u+@u>haSjDYeS5+Pe9_o0ex=^750$p}|I41dwYa)vR7 zR89JB8UimI|VOVO8ivIwGDaQqyjry*$>4vXI-t#9zn|UlPB=v9z>s z4t$Tce$8?oY5f(eUJ`i+XRc@;(Gp7Ql~8V={k0q^ayQjWX-aQJ>E`#zuFwsk%u=Dr zw}VHM|I&`9H+>30Ht2CMVwWT zzq0Tih_m?_TO3KvI>4401Ui2u@g)$v=B>~R1O`E(M&~fQq}!i-Lx*QKlhIzsY;#U^ zM)^f$;lIQpCe`I7)~DK~9L(GCkcsL!yO>FJ^mQ;Hsa3p9?ZiLC2vMcmm^TFBxOA%D zJArv0c`D7%5!%5~k9clzMTc68x$gmOj4@CGQ+#N#c%1-8w>C@CZsaPfcvdD!Zt0iGCwbu} zNfh@(%=4r;wNnDF*vLD6PEO4haj*PF|_ANE12}4Xgg-rh<6TH&uvns^V*P z#GG^ZO2oH&x7wuu#neTY0|pg6ex>5BYj2oqxatfa+3TW`cnA#lQC#iB9NdnAZ>{Ne!+KyF^o!l_h6+>Tt_r2}kc5#<0(*O_Umx>0szHFymoOd!fq3hH zUZ61u01w@qh|=qxV6$YVVv))Tgso4V574XPhAZ*M`z#1Vuo&Z^zwUm%viQq8g%ImN z02c;yHpHe%LnnaNA;k!$^(>c)g)CJ-r6T6iL{0hHRrkP5nbbg;uQ^bcszQi~n(?)s zYhxir$1 zs+H^9iTghh6LN+A2|PP4^N6zx5#-rV->qV;<>2ce1$q{um9AKTi)^Jw*)mZQ%!!bp zXSp-%~?Ba`S-K)FsBL&L@e0q=wm_7#!LF$H~f5W@%6wF`+ zbK}JJz#&1v(75+>q}TMQ$J=oCw>Mn7!W_DW?1g`-9}%+5|ExEg4ls6_Ulu&SY@_h0 zT28oTy30#=S>1s5M|L>~5348nIghm2544+_v>s(zZU!VgAkIH1&Yu_O&x!M&ibeS| z;{5C4{BCjnWpRF!IKNJue_Wi8b=-%Ilig_6=4VEz!g2dF7A7Xac5}jdt%ce5-ge!> zfx*+r_pWxwh0~f?^Y4T0UI|s&H5gW%e?TzI+icF<823~l=XaaNd&ONAG=#20;vM1+ z3nuxWnpXwKJre+fM3_bDVa={P{0}wS4Gnx(IiK=d=$w5>IZWbgUR!}h-qAc z`E;=G6+Pyex=g#YDd%j9^_z1BdscS%`!7`i#S61#0tw{7deV3pkBDR8XI#=3_k}(J z&3}y<vb(;-Yi7A~S#Mc`xdUc|`oLUeR;&WcFH)s#Y{Et(Isw=DKfrHX}T@1&Ny~)NU%&zL;R& z$|t@5ggIs$F+cRrd;Xfgu;?NEYUaP7am5m+{TPpv+JO;&oaybDStT9$?T2^cKP;t` z2%jgLE^0%tOs9FxCA|^V#Xm6zJ}LzsJI=EQ>BAHxZSI}&cgpC$TFc2Lbn0!!goZ08 z<>x^%jop~@Lh%iC6wWEN=-RsGP)_0%p*D1JHs?A{XZI~ z^hJw9X-F*{c@({FkZ#a{QCWILK<$C`9f zkIcbZ;ZRu6P=XBMk=*O!b@VacCpZVXmX>c#vfmX4 zMh@fcE)whp9qj_Yu@$#TiqIh?Ry+0jjXKUz@*>6^?eEskyQzxMvm&KA90+YO?+B(Z zC@0j40|w$<|IqOTr63z4UY(;#{6*f9hJXNk-5->}S^1!=Vvzj#D3GcDmTh3v%~Niu zSaoZbJ1^F_9FffweS(^RM9Gqeki)4~fQ`jhomXu8CZN>W*QsvBa!@;nz-%a|4rCu6 z#l~`1g0U1`$QtQP5tlAYf(2wbI)q}jt2QmKTX^9Zm|VARv@S_1V5vHum|G65SpbPW zmM&L5=9MY&>+V{}V353YZ%4|Wtv=!$Tu|p<;QYSnGl8>8TxOhuKw9^P@>o?;2ka2h zo5KMmqn2=>=lf-(_kR!$Q$zPE33FkZ;J4~v?3EM73wg^83Eg{EV)+xQ5-ilUpeKz9 zYI}3tJIe2v?cSWj0Zw$s#~#E}O8myV+qP0NZ6a~Z;u&~%q6?^Z14!G!A}q%i;S-k- z_u83L$xlzc6V0C5r78@*i@Y`pWiEoiT6SfDV`ZpYC8{EB0`-=DPN{B{0#|y&Zn{e> zqk+}jcnUqsW*RhBxWq=*rd7^oSBCGaOwR^_t+9@a23PWrs;dtryY(^LPBYwY3+$Ek z%wmQ?CQ68E7z&lRmB`#`&bUpE)71Cbk(|_VEPSWEE}StXsa7f(?zMMm@#`WvQRN5- zku+q=jjR&%a-;kWhi2~v29?y}2k1y%3p6hsnxk67_!5sYKt~@n1DDsI$mRDu9tr4C zSor_sssRNxR)1zcSs4{*Ier&VQ{37pfyVP`F^>S+h5ycr+GMF?Y?xtq&)l_{Tox_h zdF>-zGsaI-g||GmBkBk5-|l)h#Zr2V?QQqGOJdbgG@lm?pSD=q-|bhH`INse{NASL~RH0#yS!m2*ptv;hY2#0aV?=DUgTeqzkO z^P>OsiyB7=%(H=TtMV(2`puQH_S?u(dL^D*z|W7_?IEgvZyBEnSfw zqGRs9#uDxt@KVWQs({k-zn^f8(u4;kBS+xqS{irZcLFM9MZZ`zh`Y33>M&O<~7OFi29 zA2+Jo+1IanTz79Om0ge~N&QuK$Zajkb|cqM^tXtRm)J4`t_Nqp$U<-^%Ac_U>qY`e zN14*n9v;~p*K+ywXnQ!-3_E<3di6TFeTR7#@(^*K^=M1I4DLVZL+DwY^&yb46Gd1$ zAMII)Z1}yoJK{uF`u?+QYuMe>mnM{C%63*~Op$hzBR7Yey0g6O4GF#{`c?ye;+)q+ zsMkb)Q|jqEQY12J%hHP9YMYQYal4mKZu6QPZOVCtWK#0?Dz(ykTx*?Q2dRU02up5x zI7Kc_fxU`2d6KeRN!cNg-K>8?StqM=OBZrbFO|5v3wb0>DoTM&QhPPQpPPcYey*`L z`p8gJ&!L>^wz+d{56*dAai8h#l1>?ABPV+D-BoQ1s~Gk)?^8m@x+yC|cDt5x-DL;vW6Hqr&VlY7 zq(@d`j6yBfM7Mx{Wr`25JgWRrRJN3@UZs4yd57j>jL2!vfM$KRx@mpnLOeSK%T`;C z@JKTah7l(dKD!c+ru>|d@@q!Qr|~zB@<=~7Fe|MMqkgex_Rk%>0+**%DNp4IF6GAd zFZVnpB}fYFpEaaBt4f%xv>sy&hr-_}U{QXgj|%&i0b&3tPWH;ksdgB4ZR$|6q7A}(mT zUu1OR0;h7hW~=PZ0|BZeY&}ibdO_Dr`zx_gJla7|+h~=x#wtx0-#wTrkeo0UC8?{@`j}&hYtz=QPkS6sE75r3HiCQL z=^*uFpt>rEnwMB_J3x2}|A+g%1bjg^f)*wN7pj8Be228QWG(HJ&`eYiR>mW(yOg1@ z5rNDq+ekMI+*ET?{$)qnzMiz8p0sE2nzDVQ4X{CW#pYC5!g$(echiFJr#+9ypEPFH zkM5^EVJ=KvSg1a=n0U16Od2POYi{V&S97ZB%NJ?uU!{e-O8b2s5P%Cyw)yDB%CA%p zFxMA>>qo1skjxZbHRk&kD9jYs81{NgR(ekgujWExO;u9Oodg>oD>llqvNG?dXlj^Q z`I%Xnd$4+G45uEoosa~NYl0`PPuS>vm@qJ43hu+EYzBxfJ}8XAqRKH7Zo)a@lm+bf zkLp|{JS+b{ZFBpT@Dn+hg8nML@?rJq2SS3Yd;)B2`NU^n<)r`go#lh#N?&KL$`NMn z`FK?pZ~@LVHw6b`<~bn*?7TjNs~(>J#rpi@1MIBrU!MT(@riFag1wYC${!?~#(d{R zvJhhTluaC(W$^y<%}pRaaxt}s+C(qcoRu)~zEi#fpjiDFx=xhwlEyK{=7+SLtL)S-OU^@x`E1fa$k(HS08a=oI10P}l@JuuY zp8EiRiv$>0=#GJtT`;3tG4Mzg2L6swzEg9yc9&jt6sG#H9_^#I14GNc!_f2^483|8 gQ050)a2OqW}N^ diff --git a/piet-gpu/shader/gen/coarse.hlsl b/piet-gpu/shader/gen/coarse.hlsl index 3a5c1a5..a7f769f 100644 --- a/piet-gpu/shader/gen/coarse.hlsl +++ b/piet-gpu/shader/gen/coarse.hlsl @@ -175,9 +175,9 @@ struct Config static const uint3 gl_WorkGroupSize = uint3(256u, 1u, 1u); -RWByteAddressBuffer _260 : register(u0, space0); -ByteAddressBuffer _1005 : register(t1, space0); -ByteAddressBuffer _1378 : register(t2, space0); +RWByteAddressBuffer _266 : register(u0, space0); +ByteAddressBuffer _1020 : register(t1, space0); +ByteAddressBuffer _1399 : register(t2, space0); static uint3 gl_WorkGroupID; static uint3 gl_LocalInvocationID; @@ -200,8 +200,8 @@ groupshared uint sh_tile_count[256]; Alloc slice_mem(Alloc a, uint offset, uint size) { - Alloc _337 = { a.offset + offset }; - return _337; + Alloc _343 = { a.offset + offset }; + return _343; } bool touch_mem(Alloc alloc, uint offset) @@ -217,7 +217,7 @@ uint read_mem(Alloc alloc, uint offset) { return 0u; } - uint v = _260.Load(offset * 4 + 8); + uint v = _266.Load(offset * 4 + 8); return v; } @@ -230,8 +230,8 @@ Alloc new_alloc(uint offset, uint size, bool mem_ok) BinInstanceRef BinInstance_index(BinInstanceRef ref, uint index) { - BinInstanceRef _346 = { ref.offset + (index * 4u) }; - return _346; + BinInstanceRef _361 = { ref.offset + (index * 4u) }; + return _361; } BinInstance BinInstance_read(Alloc a, BinInstanceRef ref) @@ -259,8 +259,8 @@ Path Path_read(Alloc a, PathRef ref) uint raw2 = read_mem(param_4, param_5); Path s; s.bbox = uint4(raw0 & 65535u, raw0 >> uint(16), raw1 & 65535u, raw1 >> uint(16)); - TileRef _409 = { raw2 }; - s.tiles = _409; + TileRef _424 = { raw2 }; + s.tiles = _424; return s; } @@ -270,11 +270,11 @@ void write_tile_alloc(uint el_ix, Alloc a) Alloc read_tile_alloc(uint el_ix, bool mem_ok) { - uint _892; - _260.GetDimensions(_892); - _892 = (_892 - 8) / 4; + uint _907; + _266.GetDimensions(_907); + _907 = (_907 - 8) / 4; uint param = 0u; - uint param_1 = uint(int(_892) * 4); + uint param_1 = uint(int(_907) * 4); bool param_2 = mem_ok; return new_alloc(param, param_1, param_2); } @@ -288,31 +288,31 @@ Tile Tile_read(Alloc a, TileRef ref) Alloc param_2 = a; uint param_3 = ix + 1u; uint raw1 = read_mem(param_2, param_3); - TileSegRef _434 = { raw0 }; + TileSegRef _449 = { raw0 }; Tile s; - s.tile = _434; + s.tile = _449; s.backdrop = int(raw1); return s; } MallocResult malloc(uint size) { - uint _266; - _260.InterlockedAdd(0, size, _266); - uint offset = _266; - uint _273; - _260.GetDimensions(_273); - _273 = (_273 - 8) / 4; + uint _272; + _266.InterlockedAdd(0, size, _272); + uint offset = _272; + uint _279; + _266.GetDimensions(_279); + _279 = (_279 - 8) / 4; MallocResult r; - r.failed = (offset + size) > uint(int(_273) * 4); + r.failed = (offset + size) > uint(int(_279) * 4); uint param = offset; uint param_1 = size; bool param_2 = !r.failed; r.alloc = new_alloc(param, param_1, param_2); if (r.failed) { - uint _295; - _260.InterlockedMax(4, 1u, _295); + uint _301; + _266.InterlockedMax(4, 1u, _301); return r; } return r; @@ -326,7 +326,7 @@ void write_mem(Alloc alloc, uint offset, uint val) { return; } - _260.Store(offset * 4 + 8, val); + _266.Store(offset * 4 + 8, val); } void CmdJump_write(Alloc a, CmdJumpRef ref, CmdJump s) @@ -344,9 +344,9 @@ void Cmd_Jump_write(Alloc a, CmdRef ref, CmdJump s) uint param_1 = ref.offset >> uint(2); uint param_2 = 11u; write_mem(param, param_1, param_2); - CmdJumpRef _885 = { ref.offset + 4u }; + CmdJumpRef _900 = { ref.offset + 4u }; Alloc param_3 = a; - CmdJumpRef param_4 = _885; + CmdJumpRef param_4 = _900; CmdJump param_5 = s; CmdJump_write(param_3, param_4, param_5); } @@ -358,21 +358,21 @@ bool alloc_cmd(inout Alloc cmd_alloc, inout CmdRef cmd_ref, inout uint cmd_limit return true; } uint param = 1024u; - MallocResult _913 = malloc(param); - MallocResult new_cmd = _913; + MallocResult _928 = malloc(param); + MallocResult new_cmd = _928; if (new_cmd.failed) { return false; } - CmdJump _923 = { new_cmd.alloc.offset }; - CmdJump jump = _923; + CmdJump _938 = { new_cmd.alloc.offset }; + CmdJump jump = _938; Alloc param_1 = cmd_alloc; CmdRef param_2 = cmd_ref; CmdJump param_3 = jump; Cmd_Jump_write(param_1, param_2, param_3); cmd_alloc = new_cmd.alloc; - CmdRef _935 = { cmd_alloc.offset }; - cmd_ref = _935; + CmdRef _950 = { cmd_alloc.offset }; + cmd_ref = _950; cmd_limit = (cmd_alloc.offset + 1024u) - 144u; return true; } @@ -396,9 +396,9 @@ void Cmd_Fill_write(Alloc a, CmdRef ref, CmdFill s) uint param_1 = ref.offset >> uint(2); uint param_2 = 1u; write_mem(param, param_1, param_2); - CmdFillRef _742 = { ref.offset + 4u }; + CmdFillRef _757 = { ref.offset + 4u }; Alloc param_3 = a; - CmdFillRef param_4 = _742; + CmdFillRef param_4 = _757; CmdFill param_5 = s; CmdFill_write(param_3, param_4, param_5); } @@ -430,9 +430,9 @@ void Cmd_Stroke_write(Alloc a, CmdRef ref, CmdStroke s) uint param_1 = ref.offset >> uint(2); uint param_2 = 2u; write_mem(param, param_1, param_2); - CmdStrokeRef _760 = { ref.offset + 4u }; + CmdStrokeRef _775 = { ref.offset + 4u }; Alloc param_3 = a; - CmdStrokeRef param_4 = _760; + CmdStrokeRef param_4 = _775; CmdStroke param_5 = s; CmdStroke_write(param_3, param_4, param_5); } @@ -443,8 +443,8 @@ void write_fill(Alloc alloc, inout CmdRef cmd_ref, Tile tile, float linewidth) { if (tile.tile.offset != 0u) { - CmdFill _958 = { tile.tile.offset, tile.backdrop }; - CmdFill cmd_fill = _958; + CmdFill _973 = { tile.tile.offset, tile.backdrop }; + CmdFill cmd_fill = _973; Alloc param = alloc; CmdRef param_1 = cmd_ref; CmdFill param_2 = cmd_fill; @@ -461,8 +461,8 @@ void write_fill(Alloc alloc, inout CmdRef cmd_ref, Tile tile, float linewidth) } else { - CmdStroke _988 = { tile.tile.offset, 0.5f * linewidth }; - CmdStroke cmd_stroke = _988; + CmdStroke _1003 = { tile.tile.offset, 0.5f * linewidth }; + CmdStroke cmd_stroke = _1003; Alloc param_5 = alloc; CmdRef param_6 = cmd_ref; CmdStroke param_7 = cmd_stroke; @@ -486,9 +486,9 @@ void Cmd_Color_write(Alloc a, CmdRef ref, CmdColor s) uint param_1 = ref.offset >> uint(2); uint param_2 = 5u; write_mem(param, param_1, param_2); - CmdColorRef _786 = { ref.offset + 4u }; + CmdColorRef _801 = { ref.offset + 4u }; Alloc param_3 = a; - CmdColorRef param_4 = _786; + CmdColorRef param_4 = _801; CmdColor param_5 = s; CmdColor_write(param_3, param_4, param_5); } @@ -520,9 +520,9 @@ void Cmd_LinGrad_write(Alloc a, CmdRef ref, CmdLinGrad s) uint param_1 = ref.offset >> uint(2); uint param_2 = 6u; write_mem(param, param_1, param_2); - CmdLinGradRef _804 = { ref.offset + 4u }; + CmdLinGradRef _819 = { ref.offset + 4u }; Alloc param_3 = a; - CmdLinGradRef param_4 = _804; + CmdLinGradRef param_4 = _819; CmdLinGrad param_5 = s; CmdLinGrad_write(param_3, param_4, param_5); } @@ -582,9 +582,9 @@ void Cmd_RadGrad_write(Alloc a, CmdRef ref, CmdRadGrad s) uint param_1 = ref.offset >> uint(2); uint param_2 = 7u; write_mem(param, param_1, param_2); - CmdRadGradRef _822 = { ref.offset + 4u }; + CmdRadGradRef _837 = { ref.offset + 4u }; Alloc param_3 = a; - CmdRadGradRef param_4 = _822; + CmdRadGradRef param_4 = _837; CmdRadGrad param_5 = s; CmdRadGrad_write(param_3, param_4, param_5); } @@ -608,9 +608,9 @@ void Cmd_Image_write(Alloc a, CmdRef ref, CmdImage s) uint param_1 = ref.offset >> uint(2); uint param_2 = 8u; write_mem(param, param_1, param_2); - CmdImageRef _840 = { ref.offset + 4u }; + CmdImageRef _855 = { ref.offset + 4u }; Alloc param_3 = a; - CmdImageRef param_4 = _840; + CmdImageRef param_4 = _855; CmdImage param_5 = s; CmdImage_write(param_3, param_4, param_5); } @@ -638,9 +638,9 @@ void Cmd_EndClip_write(Alloc a, CmdRef ref, CmdEndClip s) uint param_1 = ref.offset >> uint(2); uint param_2 = 10u; write_mem(param, param_1, param_2); - CmdEndClipRef _866 = { ref.offset + 4u }; + CmdEndClipRef _881 = { ref.offset + 4u }; Alloc param_3 = a; - CmdEndClipRef param_4 = _866; + CmdEndClipRef param_4 = _881; CmdEndClip param_5 = s; CmdEndClip_write(param_3, param_4, param_5); } @@ -653,27 +653,35 @@ void Cmd_End_write(Alloc a, CmdRef ref) write_mem(param, param_1, param_2); } +void alloc_write(Alloc a, uint offset, Alloc alloc) +{ + Alloc param = a; + uint param_1 = offset >> uint(2); + uint param_2 = alloc.offset; + write_mem(param, param_1, param_2); +} + void comp_main() { - uint width_in_bins = ((_1005.Load(8) + 16u) - 1u) / 16u; + uint width_in_bins = ((_1020.Load(8) + 16u) - 1u) / 16u; uint bin_ix = (width_in_bins * gl_WorkGroupID.y) + gl_WorkGroupID.x; uint partition_ix = 0u; - uint n_partitions = ((_1005.Load(0) + 256u) - 1u) / 256u; + uint n_partitions = ((_1020.Load(0) + 256u) - 1u) / 256u; uint th_ix = gl_LocalInvocationID.x; uint bin_tile_x = 16u * gl_WorkGroupID.x; uint bin_tile_y = 16u * gl_WorkGroupID.y; uint tile_x = gl_LocalInvocationID.x % 16u; uint tile_y = gl_LocalInvocationID.x / 16u; - uint this_tile_ix = (((bin_tile_y + tile_y) * _1005.Load(8)) + bin_tile_x) + tile_x; - Alloc _1070; - _1070.offset = _1005.Load(24); + uint this_tile_ix = (((bin_tile_y + tile_y) * _1020.Load(8)) + bin_tile_x) + tile_x; + Alloc _1085; + _1085.offset = _1020.Load(24); Alloc param; - param.offset = _1070.offset; + param.offset = _1085.offset; uint param_1 = this_tile_ix * 1024u; uint param_2 = 1024u; Alloc cmd_alloc = slice_mem(param, param_1, param_2); - CmdRef _1079 = { cmd_alloc.offset }; - CmdRef cmd_ref = _1079; + CmdRef _1094 = { cmd_alloc.offset }; + CmdRef cmd_ref = _1094; uint cmd_limit = (cmd_ref.offset + 1024u) - 144u; uint clip_depth = 0u; uint clip_zero_depth = 0u; @@ -681,21 +689,25 @@ void comp_main() uint wr_ix = 0u; uint part_start_ix = 0u; uint ready_ix = 0u; + Alloc param_3 = cmd_alloc; + uint param_4 = 0u; + uint param_5 = 8u; + Alloc scratch_alloc = slice_mem(param_3, param_4, param_5); cmd_ref.offset += 4u; uint render_blend_depth = 0u; uint max_blend_depth = 0u; - uint drawmonoid_start = _1005.Load(44) >> uint(2); - uint drawtag_start = _1005.Load(100) >> uint(2); - uint drawdata_start = _1005.Load(104) >> uint(2); - uint drawinfo_start = _1005.Load(68) >> uint(2); - bool mem_ok = _260.Load(4) == 0u; - Alloc param_3; - Alloc param_5; - uint _1310; + uint drawmonoid_start = _1020.Load(44) >> uint(2); + uint drawtag_start = _1020.Load(100) >> uint(2); + uint drawdata_start = _1020.Load(104) >> uint(2); + uint drawinfo_start = _1020.Load(68) >> uint(2); + bool mem_ok = _266.Load(4) == 0u; + Alloc param_6; + Alloc param_8; + uint _1331; uint element_ix; - Alloc param_14; + Alloc param_17; uint tile_count; - uint _1611; + uint _1632; float linewidth; CmdLinGrad cmd_lin; CmdRadGrad cmd_rad; @@ -705,40 +717,40 @@ void comp_main() { sh_bitmaps[i][th_ix] = 0u; } - bool _1362; + bool _1383; for (;;) { if ((ready_ix == wr_ix) && (partition_ix < n_partitions)) { part_start_ix = ready_ix; uint count = 0u; - bool _1160 = th_ix < 256u; - bool _1168; - if (_1160) + bool _1181 = th_ix < 256u; + bool _1189; + if (_1181) { - _1168 = (partition_ix + th_ix) < n_partitions; + _1189 = (partition_ix + th_ix) < n_partitions; } else { - _1168 = _1160; + _1189 = _1181; } - if (_1168) + if (_1189) { - uint in_ix = (_1005.Load(20) >> uint(2)) + ((((partition_ix + th_ix) * 256u) + bin_ix) * 2u); - Alloc _1185; - _1185.offset = _1005.Load(20); - param_3.offset = _1185.offset; - uint param_4 = in_ix; - count = read_mem(param_3, param_4); - Alloc _1196; - _1196.offset = _1005.Load(20); - param_5.offset = _1196.offset; - uint param_6 = in_ix + 1u; - uint offset = read_mem(param_5, param_6); - uint param_7 = offset; - uint param_8 = count * 4u; - bool param_9 = mem_ok; - sh_part_elements[th_ix] = new_alloc(param_7, param_8, param_9); + uint in_ix = (_1020.Load(20) >> uint(2)) + ((((partition_ix + th_ix) * 256u) + bin_ix) * 2u); + Alloc _1206; + _1206.offset = _1020.Load(20); + param_6.offset = _1206.offset; + uint param_7 = in_ix; + count = read_mem(param_6, param_7); + Alloc _1217; + _1217.offset = _1020.Load(20); + param_8.offset = _1217.offset; + uint param_9 = in_ix + 1u; + uint offset = read_mem(param_8, param_9); + uint param_10 = offset; + uint param_11 = count * 4u; + bool param_12 = mem_ok; + sh_part_elements[th_ix] = new_alloc(param_10, param_11, param_12); } for (uint i_1 = 0u; i_1 < 8u; i_1++) { @@ -778,35 +790,35 @@ void comp_main() } if (part_ix > 0u) { - _1310 = sh_part_count[part_ix - 1u]; + _1331 = sh_part_count[part_ix - 1u]; } else { - _1310 = part_start_ix; + _1331 = part_start_ix; } - ix -= _1310; + ix -= _1331; Alloc bin_alloc = sh_part_elements[part_ix]; - BinInstanceRef _1329 = { bin_alloc.offset }; - BinInstanceRef inst_ref = _1329; - BinInstanceRef param_10 = inst_ref; - uint param_11 = ix; - Alloc param_12 = bin_alloc; - BinInstanceRef param_13 = BinInstance_index(param_10, param_11); - BinInstance inst = BinInstance_read(param_12, param_13); + BinInstanceRef _1350 = { bin_alloc.offset }; + BinInstanceRef inst_ref = _1350; + BinInstanceRef param_13 = inst_ref; + uint param_14 = ix; + Alloc param_15 = bin_alloc; + BinInstanceRef param_16 = BinInstance_index(param_13, param_14); + BinInstance inst = BinInstance_read(param_15, param_16); sh_elements[th_ix] = inst.element_ix; } GroupMemoryBarrierWithGroupSync(); wr_ix = min((rd_ix + 256u), ready_ix); - bool _1352 = (wr_ix - rd_ix) < 256u; - if (_1352) + bool _1373 = (wr_ix - rd_ix) < 256u; + if (_1373) { - _1362 = (wr_ix < ready_ix) || (partition_ix < n_partitions); + _1383 = (wr_ix < ready_ix) || (partition_ix < n_partitions); } else { - _1362 = _1352; + _1383 = _1373; } - if (_1362) + if (_1383) { continue; } @@ -819,7 +831,7 @@ void comp_main() if ((th_ix + rd_ix) < wr_ix) { element_ix = sh_elements[th_ix]; - tag = _1378.Load((drawtag_start + element_ix) * 4 + 0); + tag = _1399.Load((drawtag_start + element_ix) * 4 + 0); } switch (tag) { @@ -831,13 +843,13 @@ void comp_main() case 37u: { uint drawmonoid_base = drawmonoid_start + (4u * element_ix); - uint path_ix = _260.Load(drawmonoid_base * 4 + 8); - PathRef _1403 = { _1005.Load(16) + (path_ix * 12u) }; - Alloc _1406; - _1406.offset = _1005.Load(16); - param_14.offset = _1406.offset; - PathRef param_15 = _1403; - Path path = Path_read(param_14, param_15); + uint path_ix = _266.Load(drawmonoid_base * 4 + 8); + PathRef _1424 = { _1020.Load(16) + (path_ix * 12u) }; + Alloc _1427; + _1427.offset = _1020.Load(16); + param_17.offset = _1427.offset; + PathRef param_18 = _1424; + Path path = Path_read(param_17, param_18); uint stride = path.bbox.z - path.bbox.x; sh_tile_stride[th_ix] = stride; int dx = int(path.bbox.x) - int(bin_tile_x); @@ -852,13 +864,13 @@ void comp_main() tile_count = uint(x1 - x0) * uint(y1 - y0); uint base = path.tiles.offset - (((uint(dy) * stride) + uint(dx)) * 8u); sh_tile_base[th_ix] = base; - uint param_16 = path.tiles.offset; - uint param_17 = ((path.bbox.z - path.bbox.x) * (path.bbox.w - path.bbox.y)) * 8u; - bool param_18 = mem_ok; - Alloc path_alloc = new_alloc(param_16, param_17, param_18); - uint param_19 = th_ix; - Alloc param_20 = path_alloc; - write_tile_alloc(param_19, param_20); + uint param_19 = path.tiles.offset; + uint param_20 = ((path.bbox.z - path.bbox.x) * (path.bbox.w - path.bbox.y)) * 8u; + bool param_21 = mem_ok; + Alloc path_alloc = new_alloc(param_19, param_20, param_21); + uint param_22 = th_ix; + Alloc param_23 = path_alloc; + write_tile_alloc(param_22, param_23); break; } default: @@ -892,56 +904,56 @@ void comp_main() } } uint element_ix_1 = sh_elements[el_ix]; - uint tag_1 = _1378.Load((drawtag_start + element_ix_1) * 4 + 0); + uint tag_1 = _1399.Load((drawtag_start + element_ix_1) * 4 + 0); if (el_ix > 0u) { - _1611 = sh_tile_count[el_ix - 1u]; + _1632 = sh_tile_count[el_ix - 1u]; } else { - _1611 = 0u; + _1632 = 0u; } - uint seq_ix = ix_1 - _1611; + uint seq_ix = ix_1 - _1632; uint width = sh_tile_width[el_ix]; uint x = sh_tile_x0[el_ix] + (seq_ix % width); uint y = sh_tile_y0[el_ix] + (seq_ix / width); bool include_tile = false; if (mem_ok) { - uint param_21 = el_ix; - bool param_22 = mem_ok; - TileRef _1663 = { sh_tile_base[el_ix] + (((sh_tile_stride[el_ix] * y) + x) * 8u) }; - Alloc param_23 = read_tile_alloc(param_21, param_22); - TileRef param_24 = _1663; - Tile tile = Tile_read(param_23, param_24); + uint param_24 = el_ix; + bool param_25 = mem_ok; + TileRef _1684 = { sh_tile_base[el_ix] + (((sh_tile_stride[el_ix] * y) + x) * 8u) }; + Alloc param_26 = read_tile_alloc(param_24, param_25); + TileRef param_27 = _1684; + Tile tile = Tile_read(param_26, param_27); bool is_clip = (tag_1 & 1u) != 0u; bool is_blend = false; if (is_clip) { uint drawmonoid_base_1 = drawmonoid_start + (4u * element_ix_1); - uint scene_offset = _260.Load((drawmonoid_base_1 + 2u) * 4 + 8); + uint scene_offset = _266.Load((drawmonoid_base_1 + 2u) * 4 + 8); uint dd = drawdata_start + (scene_offset >> uint(2)); - uint blend = _1378.Load(dd * 4 + 0); + uint blend = _1399.Load(dd * 4 + 0); is_blend = blend != 32771u; } - bool _1699 = tile.tile.offset != 0u; - bool _1708; - if (!_1699) + bool _1720 = tile.tile.offset != 0u; + bool _1729; + if (!_1720) { - _1708 = (tile.backdrop == 0) == is_clip; + _1729 = (tile.backdrop == 0) == is_clip; } else { - _1708 = _1699; + _1729 = _1720; } - include_tile = _1708 || is_blend; + include_tile = _1729 || is_blend; } if (include_tile) { uint el_slice = el_ix / 32u; uint el_mask = 1u << (el_ix & 31u); - uint _1730; - InterlockedOr(sh_bitmaps[el_slice][(y * 16u) + x], el_mask, _1730); + uint _1751; + InterlockedOr(sh_bitmaps[el_slice][(y * 16u) + x], el_mask, _1751); } } GroupMemoryBarrierWithGroupSync(); @@ -965,178 +977,178 @@ void comp_main() uint element_ref_ix = (slice_ix * 32u) + uint(int(firstbitlow(bitmap))); uint element_ix_2 = sh_elements[element_ref_ix]; bitmap &= (bitmap - 1u); - uint drawtag = _1378.Load((drawtag_start + element_ix_2) * 4 + 0); + uint drawtag = _1399.Load((drawtag_start + element_ix_2) * 4 + 0); if (clip_zero_depth == 0u) { - uint param_25 = element_ref_ix; - bool param_26 = mem_ok; - TileRef _1807 = { sh_tile_base[element_ref_ix] + (((sh_tile_stride[element_ref_ix] * tile_y) + tile_x) * 8u) }; - Alloc param_27 = read_tile_alloc(param_25, param_26); - TileRef param_28 = _1807; - Tile tile_1 = Tile_read(param_27, param_28); + uint param_28 = element_ref_ix; + bool param_29 = mem_ok; + TileRef _1828 = { sh_tile_base[element_ref_ix] + (((sh_tile_stride[element_ref_ix] * tile_y) + tile_x) * 8u) }; + Alloc param_30 = read_tile_alloc(param_28, param_29); + TileRef param_31 = _1828; + Tile tile_1 = Tile_read(param_30, param_31); uint drawmonoid_base_2 = drawmonoid_start + (4u * element_ix_2); - uint scene_offset_1 = _260.Load((drawmonoid_base_2 + 2u) * 4 + 8); - uint info_offset = _260.Load((drawmonoid_base_2 + 3u) * 4 + 8); + uint scene_offset_1 = _266.Load((drawmonoid_base_2 + 2u) * 4 + 8); + uint info_offset = _266.Load((drawmonoid_base_2 + 3u) * 4 + 8); uint dd_1 = drawdata_start + (scene_offset_1 >> uint(2)); uint di = drawinfo_start + (info_offset >> uint(2)); switch (drawtag) { case 68u: { - linewidth = asfloat(_260.Load(di * 4 + 8)); - Alloc param_29 = cmd_alloc; - CmdRef param_30 = cmd_ref; - uint param_31 = cmd_limit; - bool _1855 = alloc_cmd(param_29, param_30, param_31); - cmd_alloc = param_29; - cmd_ref = param_30; - cmd_limit = param_31; - if (!_1855) + linewidth = asfloat(_266.Load(di * 4 + 8)); + Alloc param_32 = cmd_alloc; + CmdRef param_33 = cmd_ref; + uint param_34 = cmd_limit; + bool _1876 = alloc_cmd(param_32, param_33, param_34); + cmd_alloc = param_32; + cmd_ref = param_33; + cmd_limit = param_34; + if (!_1876) { break; } - Alloc param_32 = cmd_alloc; - CmdRef param_33 = cmd_ref; - Tile param_34 = tile_1; - float param_35 = linewidth; - write_fill(param_32, param_33, param_34, param_35); - cmd_ref = param_33; - uint rgba = _1378.Load(dd_1 * 4 + 0); - CmdColor _1878 = { rgba }; - Alloc param_36 = cmd_alloc; - CmdRef param_37 = cmd_ref; - CmdColor param_38 = _1878; - Cmd_Color_write(param_36, param_37, param_38); + Alloc param_35 = cmd_alloc; + CmdRef param_36 = cmd_ref; + Tile param_37 = tile_1; + float param_38 = linewidth; + write_fill(param_35, param_36, param_37, param_38); + cmd_ref = param_36; + uint rgba = _1399.Load(dd_1 * 4 + 0); + CmdColor _1899 = { rgba }; + Alloc param_39 = cmd_alloc; + CmdRef param_40 = cmd_ref; + CmdColor param_41 = _1899; + Cmd_Color_write(param_39, param_40, param_41); cmd_ref.offset += 8u; break; } case 276u: { - Alloc param_39 = cmd_alloc; - CmdRef param_40 = cmd_ref; - uint param_41 = cmd_limit; - bool _1896 = alloc_cmd(param_39, param_40, param_41); - cmd_alloc = param_39; - cmd_ref = param_40; - cmd_limit = param_41; - if (!_1896) + Alloc param_42 = cmd_alloc; + CmdRef param_43 = cmd_ref; + uint param_44 = cmd_limit; + bool _1917 = alloc_cmd(param_42, param_43, param_44); + cmd_alloc = param_42; + cmd_ref = param_43; + cmd_limit = param_44; + if (!_1917) { break; } - linewidth = asfloat(_260.Load(di * 4 + 8)); - Alloc param_42 = cmd_alloc; - CmdRef param_43 = cmd_ref; - Tile param_44 = tile_1; - float param_45 = linewidth; - write_fill(param_42, param_43, param_44, param_45); - cmd_ref = param_43; - cmd_lin.index = _1378.Load(dd_1 * 4 + 0); - cmd_lin.line_x = asfloat(_260.Load((di + 1u) * 4 + 8)); - cmd_lin.line_y = asfloat(_260.Load((di + 2u) * 4 + 8)); - cmd_lin.line_c = asfloat(_260.Load((di + 3u) * 4 + 8)); - Alloc param_46 = cmd_alloc; - CmdRef param_47 = cmd_ref; - CmdLinGrad param_48 = cmd_lin; - Cmd_LinGrad_write(param_46, param_47, param_48); + linewidth = asfloat(_266.Load(di * 4 + 8)); + Alloc param_45 = cmd_alloc; + CmdRef param_46 = cmd_ref; + Tile param_47 = tile_1; + float param_48 = linewidth; + write_fill(param_45, param_46, param_47, param_48); + cmd_ref = param_46; + cmd_lin.index = _1399.Load(dd_1 * 4 + 0); + cmd_lin.line_x = asfloat(_266.Load((di + 1u) * 4 + 8)); + cmd_lin.line_y = asfloat(_266.Load((di + 2u) * 4 + 8)); + cmd_lin.line_c = asfloat(_266.Load((di + 3u) * 4 + 8)); + Alloc param_49 = cmd_alloc; + CmdRef param_50 = cmd_ref; + CmdLinGrad param_51 = cmd_lin; + Cmd_LinGrad_write(param_49, param_50, param_51); cmd_ref.offset += 20u; break; } case 732u: { - Alloc param_49 = cmd_alloc; - CmdRef param_50 = cmd_ref; - uint param_51 = cmd_limit; - bool _1960 = alloc_cmd(param_49, param_50, param_51); - cmd_alloc = param_49; - cmd_ref = param_50; - cmd_limit = param_51; - if (!_1960) + Alloc param_52 = cmd_alloc; + CmdRef param_53 = cmd_ref; + uint param_54 = cmd_limit; + bool _1981 = alloc_cmd(param_52, param_53, param_54); + cmd_alloc = param_52; + cmd_ref = param_53; + cmd_limit = param_54; + if (!_1981) { break; } - linewidth = asfloat(_260.Load(di * 4 + 8)); - Alloc param_52 = cmd_alloc; - CmdRef param_53 = cmd_ref; - Tile param_54 = tile_1; - float param_55 = linewidth; - write_fill(param_52, param_53, param_54, param_55); - cmd_ref = param_53; - cmd_rad.index = _1378.Load(dd_1 * 4 + 0); - cmd_rad.mat = asfloat(uint4(_260.Load((di + 1u) * 4 + 8), _260.Load((di + 2u) * 4 + 8), _260.Load((di + 3u) * 4 + 8), _260.Load((di + 4u) * 4 + 8))); - cmd_rad.xlat = asfloat(uint2(_260.Load((di + 5u) * 4 + 8), _260.Load((di + 6u) * 4 + 8))); - cmd_rad.c1 = asfloat(uint2(_260.Load((di + 7u) * 4 + 8), _260.Load((di + 8u) * 4 + 8))); - cmd_rad.ra = asfloat(_260.Load((di + 9u) * 4 + 8)); - cmd_rad.roff = asfloat(_260.Load((di + 10u) * 4 + 8)); - Alloc param_56 = cmd_alloc; - CmdRef param_57 = cmd_ref; - CmdRadGrad param_58 = cmd_rad; - Cmd_RadGrad_write(param_56, param_57, param_58); + linewidth = asfloat(_266.Load(di * 4 + 8)); + Alloc param_55 = cmd_alloc; + CmdRef param_56 = cmd_ref; + Tile param_57 = tile_1; + float param_58 = linewidth; + write_fill(param_55, param_56, param_57, param_58); + cmd_ref = param_56; + cmd_rad.index = _1399.Load(dd_1 * 4 + 0); + cmd_rad.mat = asfloat(uint4(_266.Load((di + 1u) * 4 + 8), _266.Load((di + 2u) * 4 + 8), _266.Load((di + 3u) * 4 + 8), _266.Load((di + 4u) * 4 + 8))); + cmd_rad.xlat = asfloat(uint2(_266.Load((di + 5u) * 4 + 8), _266.Load((di + 6u) * 4 + 8))); + cmd_rad.c1 = asfloat(uint2(_266.Load((di + 7u) * 4 + 8), _266.Load((di + 8u) * 4 + 8))); + cmd_rad.ra = asfloat(_266.Load((di + 9u) * 4 + 8)); + cmd_rad.roff = asfloat(_266.Load((di + 10u) * 4 + 8)); + Alloc param_59 = cmd_alloc; + CmdRef param_60 = cmd_ref; + CmdRadGrad param_61 = cmd_rad; + Cmd_RadGrad_write(param_59, param_60, param_61); cmd_ref.offset += 48u; break; } case 72u: { - linewidth = asfloat(_260.Load(di * 4 + 8)); - Alloc param_59 = cmd_alloc; - CmdRef param_60 = cmd_ref; - uint param_61 = cmd_limit; - bool _2066 = alloc_cmd(param_59, param_60, param_61); - cmd_alloc = param_59; - cmd_ref = param_60; - cmd_limit = param_61; - if (!_2066) + linewidth = asfloat(_266.Load(di * 4 + 8)); + Alloc param_62 = cmd_alloc; + CmdRef param_63 = cmd_ref; + uint param_64 = cmd_limit; + bool _2087 = alloc_cmd(param_62, param_63, param_64); + cmd_alloc = param_62; + cmd_ref = param_63; + cmd_limit = param_64; + if (!_2087) { break; } - Alloc param_62 = cmd_alloc; - CmdRef param_63 = cmd_ref; - Tile param_64 = tile_1; - float param_65 = linewidth; - write_fill(param_62, param_63, param_64, param_65); - cmd_ref = param_63; - uint index = _1378.Load(dd_1 * 4 + 0); - uint raw1 = _1378.Load((dd_1 + 1u) * 4 + 0); + Alloc param_65 = cmd_alloc; + CmdRef param_66 = cmd_ref; + Tile param_67 = tile_1; + float param_68 = linewidth; + write_fill(param_65, param_66, param_67, param_68); + cmd_ref = param_66; + uint index = _1399.Load(dd_1 * 4 + 0); + uint raw1 = _1399.Load((dd_1 + 1u) * 4 + 0); int2 offset_1 = int2(int(raw1 << uint(16)) >> 16, int(raw1) >> 16); - CmdImage _2105 = { index, offset_1 }; - Alloc param_66 = cmd_alloc; - CmdRef param_67 = cmd_ref; - CmdImage param_68 = _2105; - Cmd_Image_write(param_66, param_67, param_68); + CmdImage _2126 = { index, offset_1 }; + Alloc param_69 = cmd_alloc; + CmdRef param_70 = cmd_ref; + CmdImage param_71 = _2126; + Cmd_Image_write(param_69, param_70, param_71); cmd_ref.offset += 12u; break; } case 5u: { - bool _2119 = tile_1.tile.offset == 0u; - bool _2125; - if (_2119) + bool _2140 = tile_1.tile.offset == 0u; + bool _2146; + if (_2140) { - _2125 = tile_1.backdrop == 0; + _2146 = tile_1.backdrop == 0; } else { - _2125 = _2119; + _2146 = _2140; } - if (_2125) + if (_2146) { clip_zero_depth = clip_depth + 1u; } else { - Alloc param_69 = cmd_alloc; - CmdRef param_70 = cmd_ref; - uint param_71 = cmd_limit; - bool _2137 = alloc_cmd(param_69, param_70, param_71); - cmd_alloc = param_69; - cmd_ref = param_70; - cmd_limit = param_71; - if (!_2137) + Alloc param_72 = cmd_alloc; + CmdRef param_73 = cmd_ref; + uint param_74 = cmd_limit; + bool _2158 = alloc_cmd(param_72, param_73, param_74); + cmd_alloc = param_72; + cmd_ref = param_73; + cmd_limit = param_74; + if (!_2158) { break; } - Alloc param_72 = cmd_alloc; - CmdRef param_73 = cmd_ref; - Cmd_BeginClip_write(param_72, param_73); + Alloc param_75 = cmd_alloc; + CmdRef param_76 = cmd_ref; + Cmd_BeginClip_write(param_75, param_76); cmd_ref.offset += 4u; render_blend_depth++; max_blend_depth = max(max_blend_depth, render_blend_depth); @@ -1147,29 +1159,29 @@ void comp_main() case 37u: { clip_depth--; - Alloc param_74 = cmd_alloc; - CmdRef param_75 = cmd_ref; - uint param_76 = cmd_limit; - bool _2170 = alloc_cmd(param_74, param_75, param_76); - cmd_alloc = param_74; - cmd_ref = param_75; - cmd_limit = param_76; - if (!_2170) + Alloc param_77 = cmd_alloc; + CmdRef param_78 = cmd_ref; + uint param_79 = cmd_limit; + bool _2191 = alloc_cmd(param_77, param_78, param_79); + cmd_alloc = param_77; + cmd_ref = param_78; + cmd_limit = param_79; + if (!_2191) { break; } - Alloc param_77 = cmd_alloc; - CmdRef param_78 = cmd_ref; - Tile param_79 = tile_1; - float param_80 = -1.0f; - write_fill(param_77, param_78, param_79, param_80); - cmd_ref = param_78; - uint blend_1 = _1378.Load(dd_1 * 4 + 0); - CmdEndClip _2193 = { blend_1 }; - Alloc param_81 = cmd_alloc; - CmdRef param_82 = cmd_ref; - CmdEndClip param_83 = _2193; - Cmd_EndClip_write(param_81, param_82, param_83); + Alloc param_80 = cmd_alloc; + CmdRef param_81 = cmd_ref; + Tile param_82 = tile_1; + float param_83 = -1.0f; + write_fill(param_80, param_81, param_82, param_83); + cmd_ref = param_81; + uint blend_1 = _1399.Load(dd_1 * 4 + 0); + CmdEndClip _2214 = { blend_1 }; + Alloc param_84 = cmd_alloc; + CmdRef param_85 = cmd_ref; + CmdEndClip param_86 = _2214; + Cmd_EndClip_write(param_84, param_85, param_86); cmd_ref.offset += 8u; render_blend_depth--; break; @@ -1204,23 +1216,31 @@ void comp_main() break; } } - bool _2242 = (bin_tile_x + tile_x) < _1005.Load(8); - bool _2251; - if (_2242) + bool _2263 = (bin_tile_x + tile_x) < _1020.Load(8); + bool _2272; + if (_2263) { - _2251 = (bin_tile_y + tile_y) < _1005.Load(12); + _2272 = (bin_tile_y + tile_y) < _1020.Load(12); } else { - _2251 = _2242; + _2272 = _2263; } - if (_2251) + if (_2272) { - Alloc param_84 = cmd_alloc; - CmdRef param_85 = cmd_ref; - Cmd_End_write(param_84, param_85); + Alloc param_87 = cmd_alloc; + CmdRef param_88 = cmd_ref; + Cmd_End_write(param_87, param_88); if (max_blend_depth > 4u) { + uint scratch_size = (((max_blend_depth * 16u) * 16u) * 1u) * 4u; + uint param_89 = scratch_size; + MallocResult _2293 = malloc(param_89); + MallocResult scratch = _2293; + Alloc param_90 = scratch_alloc; + uint param_91 = scratch_alloc.offset; + Alloc param_92 = scratch.alloc; + alloc_write(param_90, param_91, param_92); } } } diff --git a/piet-gpu/shader/gen/coarse.msl b/piet-gpu/shader/gen/coarse.msl index 94b8738..d84add1 100644 --- a/piet-gpu/shader/gen/coarse.msl +++ b/piet-gpu/shader/gen/coarse.msl @@ -226,7 +226,7 @@ bool touch_mem(thread const Alloc& alloc, thread const uint& offset) } static inline __attribute__((always_inline)) -uint read_mem(thread const Alloc& alloc, thread const uint& offset, device Memory& v_260, constant uint& v_260BufferSize) +uint read_mem(thread const Alloc& alloc, thread const uint& offset, device Memory& v_266, constant uint& v_266BufferSize) { Alloc param = alloc; uint param_1 = offset; @@ -234,7 +234,7 @@ uint read_mem(thread const Alloc& alloc, thread const uint& offset, device Memor { return 0u; } - uint v = v_260.memory[offset]; + uint v = v_266.memory[offset]; return v; } @@ -253,30 +253,30 @@ BinInstanceRef BinInstance_index(thread const BinInstanceRef& ref, thread const } static inline __attribute__((always_inline)) -BinInstance BinInstance_read(thread const Alloc& a, thread const BinInstanceRef& ref, device Memory& v_260, constant uint& v_260BufferSize) +BinInstance BinInstance_read(thread const Alloc& a, thread const BinInstanceRef& ref, device Memory& v_266, constant uint& v_266BufferSize) { uint ix = ref.offset >> uint(2); Alloc param = a; uint param_1 = ix + 0u; - uint raw0 = read_mem(param, param_1, v_260, v_260BufferSize); + uint raw0 = read_mem(param, param_1, v_266, v_266BufferSize); BinInstance s; s.element_ix = raw0; return s; } static inline __attribute__((always_inline)) -Path Path_read(thread const Alloc& a, thread const PathRef& ref, device Memory& v_260, constant uint& v_260BufferSize) +Path Path_read(thread const Alloc& a, thread const PathRef& ref, device Memory& v_266, constant uint& v_266BufferSize) { uint ix = ref.offset >> uint(2); Alloc param = a; uint param_1 = ix + 0u; - uint raw0 = read_mem(param, param_1, v_260, v_260BufferSize); + uint raw0 = read_mem(param, param_1, v_266, v_266BufferSize); Alloc param_2 = a; uint param_3 = ix + 1u; - uint raw1 = read_mem(param_2, param_3, v_260, v_260BufferSize); + uint raw1 = read_mem(param_2, param_3, v_266, v_266BufferSize); Alloc param_4 = a; uint param_5 = ix + 2u; - uint raw2 = read_mem(param_4, param_5, v_260, v_260BufferSize); + uint raw2 = read_mem(param_4, param_5, v_266, v_266BufferSize); Path s; s.bbox = uint4(raw0 & 65535u, raw0 >> uint(16), raw1 & 65535u, raw1 >> uint(16)); s.tiles = TileRef{ raw2 }; @@ -289,24 +289,24 @@ void write_tile_alloc(thread const uint& el_ix, thread const Alloc& a) } static inline __attribute__((always_inline)) -Alloc read_tile_alloc(thread const uint& el_ix, thread const bool& mem_ok, device Memory& v_260, constant uint& v_260BufferSize) +Alloc read_tile_alloc(thread const uint& el_ix, thread const bool& mem_ok, device Memory& v_266, constant uint& v_266BufferSize) { uint param = 0u; - uint param_1 = uint(int((v_260BufferSize - 8) / 4) * 4); + uint param_1 = uint(int((v_266BufferSize - 8) / 4) * 4); bool param_2 = mem_ok; return new_alloc(param, param_1, param_2); } static inline __attribute__((always_inline)) -Tile Tile_read(thread const Alloc& a, thread const TileRef& ref, device Memory& v_260, constant uint& v_260BufferSize) +Tile Tile_read(thread const Alloc& a, thread const TileRef& ref, device Memory& v_266, constant uint& v_266BufferSize) { uint ix = ref.offset >> uint(2); Alloc param = a; uint param_1 = ix + 0u; - uint raw0 = read_mem(param, param_1, v_260, v_260BufferSize); + uint raw0 = read_mem(param, param_1, v_266, v_266BufferSize); Alloc param_2 = a; uint param_3 = ix + 1u; - uint raw1 = read_mem(param_2, param_3, v_260, v_260BufferSize); + uint raw1 = read_mem(param_2, param_3, v_266, v_266BufferSize); Tile s; s.tile = TileSegRef{ raw0 }; s.backdrop = int(raw1); @@ -314,26 +314,26 @@ Tile Tile_read(thread const Alloc& a, thread const TileRef& ref, device Memory& } static inline __attribute__((always_inline)) -MallocResult malloc(thread const uint& size, device Memory& v_260, constant uint& v_260BufferSize) +MallocResult malloc(thread const uint& size, device Memory& v_266, constant uint& v_266BufferSize) { - uint _266 = atomic_fetch_add_explicit((device atomic_uint*)&v_260.mem_offset, size, memory_order_relaxed); - uint offset = _266; + uint _272 = atomic_fetch_add_explicit((device atomic_uint*)&v_266.mem_offset, size, memory_order_relaxed); + uint offset = _272; MallocResult r; - r.failed = (offset + size) > uint(int((v_260BufferSize - 8) / 4) * 4); + r.failed = (offset + size) > uint(int((v_266BufferSize - 8) / 4) * 4); uint param = offset; uint param_1 = size; bool param_2 = !r.failed; r.alloc = new_alloc(param, param_1, param_2); if (r.failed) { - uint _295 = atomic_fetch_max_explicit((device atomic_uint*)&v_260.mem_error, 1u, memory_order_relaxed); + uint _301 = atomic_fetch_max_explicit((device atomic_uint*)&v_266.mem_error, 1u, memory_order_relaxed); return r; } return r; } static inline __attribute__((always_inline)) -void write_mem(thread const Alloc& alloc, thread const uint& offset, thread const uint& val, device Memory& v_260, constant uint& v_260BufferSize) +void write_mem(thread const Alloc& alloc, thread const uint& offset, thread const uint& val, device Memory& v_266, constant uint& v_266BufferSize) { Alloc param = alloc; uint param_1 = offset; @@ -341,42 +341,42 @@ void write_mem(thread const Alloc& alloc, thread const uint& offset, thread cons { return; } - v_260.memory[offset] = val; + v_266.memory[offset] = val; } static inline __attribute__((always_inline)) -void CmdJump_write(thread const Alloc& a, thread const CmdJumpRef& ref, thread const CmdJump& s, device Memory& v_260, constant uint& v_260BufferSize) +void CmdJump_write(thread const Alloc& a, thread const CmdJumpRef& ref, thread const CmdJump& s, device Memory& v_266, constant uint& v_266BufferSize) { uint ix = ref.offset >> uint(2); Alloc param = a; uint param_1 = ix + 0u; uint param_2 = s.new_ref; - write_mem(param, param_1, param_2, v_260, v_260BufferSize); + write_mem(param, param_1, param_2, v_266, v_266BufferSize); } static inline __attribute__((always_inline)) -void Cmd_Jump_write(thread const Alloc& a, thread const CmdRef& ref, thread const CmdJump& s, device Memory& v_260, constant uint& v_260BufferSize) +void Cmd_Jump_write(thread const Alloc& a, thread const CmdRef& ref, thread const CmdJump& s, device Memory& v_266, constant uint& v_266BufferSize) { Alloc param = a; uint param_1 = ref.offset >> uint(2); uint param_2 = 11u; - write_mem(param, param_1, param_2, v_260, v_260BufferSize); + write_mem(param, param_1, param_2, v_266, v_266BufferSize); Alloc param_3 = a; CmdJumpRef param_4 = CmdJumpRef{ ref.offset + 4u }; CmdJump param_5 = s; - CmdJump_write(param_3, param_4, param_5, v_260, v_260BufferSize); + CmdJump_write(param_3, param_4, param_5, v_266, v_266BufferSize); } static inline __attribute__((always_inline)) -bool alloc_cmd(thread Alloc& cmd_alloc, thread CmdRef& cmd_ref, thread uint& cmd_limit, device Memory& v_260, constant uint& v_260BufferSize) +bool alloc_cmd(thread Alloc& cmd_alloc, thread CmdRef& cmd_ref, thread uint& cmd_limit, device Memory& v_266, constant uint& v_266BufferSize) { if (cmd_ref.offset < cmd_limit) { return true; } uint param = 1024u; - MallocResult _913 = malloc(param, v_260, v_260BufferSize); - MallocResult new_cmd = _913; + MallocResult _928 = malloc(param, v_266, v_266BufferSize); + MallocResult new_cmd = _928; if (new_cmd.failed) { return false; @@ -385,7 +385,7 @@ bool alloc_cmd(thread Alloc& cmd_alloc, thread CmdRef& cmd_ref, thread uint& cmd Alloc param_1 = cmd_alloc; CmdRef param_2 = cmd_ref; CmdJump param_3 = jump; - Cmd_Jump_write(param_1, param_2, param_3, v_260, v_260BufferSize); + Cmd_Jump_write(param_1, param_2, param_3, v_266, v_266BufferSize); cmd_alloc = new_cmd.alloc; cmd_ref = CmdRef{ cmd_alloc.offset }; cmd_limit = (cmd_alloc.offset + 1024u) - 144u; @@ -393,70 +393,70 @@ bool alloc_cmd(thread Alloc& cmd_alloc, thread CmdRef& cmd_ref, thread uint& cmd } static inline __attribute__((always_inline)) -void CmdFill_write(thread const Alloc& a, thread const CmdFillRef& ref, thread const CmdFill& s, device Memory& v_260, constant uint& v_260BufferSize) +void CmdFill_write(thread const Alloc& a, thread const CmdFillRef& ref, thread const CmdFill& s, device Memory& v_266, constant uint& v_266BufferSize) { uint ix = ref.offset >> uint(2); Alloc param = a; uint param_1 = ix + 0u; uint param_2 = s.tile_ref; - write_mem(param, param_1, param_2, v_260, v_260BufferSize); + write_mem(param, param_1, param_2, v_266, v_266BufferSize); Alloc param_3 = a; uint param_4 = ix + 1u; uint param_5 = uint(s.backdrop); - write_mem(param_3, param_4, param_5, v_260, v_260BufferSize); + write_mem(param_3, param_4, param_5, v_266, v_266BufferSize); } static inline __attribute__((always_inline)) -void Cmd_Fill_write(thread const Alloc& a, thread const CmdRef& ref, thread const CmdFill& s, device Memory& v_260, constant uint& v_260BufferSize) +void Cmd_Fill_write(thread const Alloc& a, thread const CmdRef& ref, thread const CmdFill& s, device Memory& v_266, constant uint& v_266BufferSize) { Alloc param = a; uint param_1 = ref.offset >> uint(2); uint param_2 = 1u; - write_mem(param, param_1, param_2, v_260, v_260BufferSize); + write_mem(param, param_1, param_2, v_266, v_266BufferSize); Alloc param_3 = a; CmdFillRef param_4 = CmdFillRef{ ref.offset + 4u }; CmdFill param_5 = s; - CmdFill_write(param_3, param_4, param_5, v_260, v_260BufferSize); + CmdFill_write(param_3, param_4, param_5, v_266, v_266BufferSize); } static inline __attribute__((always_inline)) -void Cmd_Solid_write(thread const Alloc& a, thread const CmdRef& ref, device Memory& v_260, constant uint& v_260BufferSize) +void Cmd_Solid_write(thread const Alloc& a, thread const CmdRef& ref, device Memory& v_266, constant uint& v_266BufferSize) { Alloc param = a; uint param_1 = ref.offset >> uint(2); uint param_2 = 3u; - write_mem(param, param_1, param_2, v_260, v_260BufferSize); + write_mem(param, param_1, param_2, v_266, v_266BufferSize); } static inline __attribute__((always_inline)) -void CmdStroke_write(thread const Alloc& a, thread const CmdStrokeRef& ref, thread const CmdStroke& s, device Memory& v_260, constant uint& v_260BufferSize) +void CmdStroke_write(thread const Alloc& a, thread const CmdStrokeRef& ref, thread const CmdStroke& s, device Memory& v_266, constant uint& v_266BufferSize) { uint ix = ref.offset >> uint(2); Alloc param = a; uint param_1 = ix + 0u; uint param_2 = s.tile_ref; - write_mem(param, param_1, param_2, v_260, v_260BufferSize); + write_mem(param, param_1, param_2, v_266, v_266BufferSize); Alloc param_3 = a; uint param_4 = ix + 1u; uint param_5 = as_type(s.half_width); - write_mem(param_3, param_4, param_5, v_260, v_260BufferSize); + write_mem(param_3, param_4, param_5, v_266, v_266BufferSize); } static inline __attribute__((always_inline)) -void Cmd_Stroke_write(thread const Alloc& a, thread const CmdRef& ref, thread const CmdStroke& s, device Memory& v_260, constant uint& v_260BufferSize) +void Cmd_Stroke_write(thread const Alloc& a, thread const CmdRef& ref, thread const CmdStroke& s, device Memory& v_266, constant uint& v_266BufferSize) { Alloc param = a; uint param_1 = ref.offset >> uint(2); uint param_2 = 2u; - write_mem(param, param_1, param_2, v_260, v_260BufferSize); + write_mem(param, param_1, param_2, v_266, v_266BufferSize); Alloc param_3 = a; CmdStrokeRef param_4 = CmdStrokeRef{ ref.offset + 4u }; CmdStroke param_5 = s; - CmdStroke_write(param_3, param_4, param_5, v_260, v_260BufferSize); + CmdStroke_write(param_3, param_4, param_5, v_266, v_266BufferSize); } static inline __attribute__((always_inline)) -void write_fill(thread const Alloc& alloc, thread CmdRef& cmd_ref, thread const Tile& tile, thread const float& linewidth, device Memory& v_260, constant uint& v_260BufferSize) +void write_fill(thread const Alloc& alloc, thread CmdRef& cmd_ref, thread const Tile& tile, thread const float& linewidth, device Memory& v_266, constant uint& v_266BufferSize) { if (linewidth < 0.0) { @@ -466,14 +466,14 @@ void write_fill(thread const Alloc& alloc, thread CmdRef& cmd_ref, thread const Alloc param = alloc; CmdRef param_1 = cmd_ref; CmdFill param_2 = cmd_fill; - Cmd_Fill_write(param, param_1, param_2, v_260, v_260BufferSize); + Cmd_Fill_write(param, param_1, param_2, v_266, v_266BufferSize); cmd_ref.offset += 12u; } else { Alloc param_3 = alloc; CmdRef param_4 = cmd_ref; - Cmd_Solid_write(param_3, param_4, v_260, v_260BufferSize); + Cmd_Solid_write(param_3, param_4, v_266, v_266BufferSize); cmd_ref.offset += 4u; } } @@ -483,201 +483,210 @@ void write_fill(thread const Alloc& alloc, thread CmdRef& cmd_ref, thread const Alloc param_5 = alloc; CmdRef param_6 = cmd_ref; CmdStroke param_7 = cmd_stroke; - Cmd_Stroke_write(param_5, param_6, param_7, v_260, v_260BufferSize); + Cmd_Stroke_write(param_5, param_6, param_7, v_266, v_266BufferSize); cmd_ref.offset += 12u; } } static inline __attribute__((always_inline)) -void CmdColor_write(thread const Alloc& a, thread const CmdColorRef& ref, thread const CmdColor& s, device Memory& v_260, constant uint& v_260BufferSize) +void CmdColor_write(thread const Alloc& a, thread const CmdColorRef& ref, thread const CmdColor& s, device Memory& v_266, constant uint& v_266BufferSize) { uint ix = ref.offset >> uint(2); Alloc param = a; uint param_1 = ix + 0u; uint param_2 = s.rgba_color; - write_mem(param, param_1, param_2, v_260, v_260BufferSize); + write_mem(param, param_1, param_2, v_266, v_266BufferSize); } static inline __attribute__((always_inline)) -void Cmd_Color_write(thread const Alloc& a, thread const CmdRef& ref, thread const CmdColor& s, device Memory& v_260, constant uint& v_260BufferSize) +void Cmd_Color_write(thread const Alloc& a, thread const CmdRef& ref, thread const CmdColor& s, device Memory& v_266, constant uint& v_266BufferSize) { Alloc param = a; uint param_1 = ref.offset >> uint(2); uint param_2 = 5u; - write_mem(param, param_1, param_2, v_260, v_260BufferSize); + write_mem(param, param_1, param_2, v_266, v_266BufferSize); Alloc param_3 = a; CmdColorRef param_4 = CmdColorRef{ ref.offset + 4u }; CmdColor param_5 = s; - CmdColor_write(param_3, param_4, param_5, v_260, v_260BufferSize); + CmdColor_write(param_3, param_4, param_5, v_266, v_266BufferSize); } static inline __attribute__((always_inline)) -void CmdLinGrad_write(thread const Alloc& a, thread const CmdLinGradRef& ref, thread const CmdLinGrad& s, device Memory& v_260, constant uint& v_260BufferSize) +void CmdLinGrad_write(thread const Alloc& a, thread const CmdLinGradRef& ref, thread const CmdLinGrad& s, device Memory& v_266, constant uint& v_266BufferSize) { uint ix = ref.offset >> uint(2); Alloc param = a; uint param_1 = ix + 0u; uint param_2 = s.index; - write_mem(param, param_1, param_2, v_260, v_260BufferSize); + write_mem(param, param_1, param_2, v_266, v_266BufferSize); Alloc param_3 = a; uint param_4 = ix + 1u; uint param_5 = as_type(s.line_x); - write_mem(param_3, param_4, param_5, v_260, v_260BufferSize); + write_mem(param_3, param_4, param_5, v_266, v_266BufferSize); Alloc param_6 = a; uint param_7 = ix + 2u; uint param_8 = as_type(s.line_y); - write_mem(param_6, param_7, param_8, v_260, v_260BufferSize); + write_mem(param_6, param_7, param_8, v_266, v_266BufferSize); Alloc param_9 = a; uint param_10 = ix + 3u; uint param_11 = as_type(s.line_c); - write_mem(param_9, param_10, param_11, v_260, v_260BufferSize); + write_mem(param_9, param_10, param_11, v_266, v_266BufferSize); } static inline __attribute__((always_inline)) -void Cmd_LinGrad_write(thread const Alloc& a, thread const CmdRef& ref, thread const CmdLinGrad& s, device Memory& v_260, constant uint& v_260BufferSize) +void Cmd_LinGrad_write(thread const Alloc& a, thread const CmdRef& ref, thread const CmdLinGrad& s, device Memory& v_266, constant uint& v_266BufferSize) { Alloc param = a; uint param_1 = ref.offset >> uint(2); uint param_2 = 6u; - write_mem(param, param_1, param_2, v_260, v_260BufferSize); + write_mem(param, param_1, param_2, v_266, v_266BufferSize); Alloc param_3 = a; CmdLinGradRef param_4 = CmdLinGradRef{ ref.offset + 4u }; CmdLinGrad param_5 = s; - CmdLinGrad_write(param_3, param_4, param_5, v_260, v_260BufferSize); + CmdLinGrad_write(param_3, param_4, param_5, v_266, v_266BufferSize); } static inline __attribute__((always_inline)) -void CmdRadGrad_write(thread const Alloc& a, thread const CmdRadGradRef& ref, thread const CmdRadGrad& s, device Memory& v_260, constant uint& v_260BufferSize) +void CmdRadGrad_write(thread const Alloc& a, thread const CmdRadGradRef& ref, thread const CmdRadGrad& s, device Memory& v_266, constant uint& v_266BufferSize) { uint ix = ref.offset >> uint(2); Alloc param = a; uint param_1 = ix + 0u; uint param_2 = s.index; - write_mem(param, param_1, param_2, v_260, v_260BufferSize); + write_mem(param, param_1, param_2, v_266, v_266BufferSize); Alloc param_3 = a; uint param_4 = ix + 1u; uint param_5 = as_type(s.mat.x); - write_mem(param_3, param_4, param_5, v_260, v_260BufferSize); + write_mem(param_3, param_4, param_5, v_266, v_266BufferSize); Alloc param_6 = a; uint param_7 = ix + 2u; uint param_8 = as_type(s.mat.y); - write_mem(param_6, param_7, param_8, v_260, v_260BufferSize); + write_mem(param_6, param_7, param_8, v_266, v_266BufferSize); Alloc param_9 = a; uint param_10 = ix + 3u; uint param_11 = as_type(s.mat.z); - write_mem(param_9, param_10, param_11, v_260, v_260BufferSize); + write_mem(param_9, param_10, param_11, v_266, v_266BufferSize); Alloc param_12 = a; uint param_13 = ix + 4u; uint param_14 = as_type(s.mat.w); - write_mem(param_12, param_13, param_14, v_260, v_260BufferSize); + write_mem(param_12, param_13, param_14, v_266, v_266BufferSize); Alloc param_15 = a; uint param_16 = ix + 5u; uint param_17 = as_type(s.xlat.x); - write_mem(param_15, param_16, param_17, v_260, v_260BufferSize); + write_mem(param_15, param_16, param_17, v_266, v_266BufferSize); Alloc param_18 = a; uint param_19 = ix + 6u; uint param_20 = as_type(s.xlat.y); - write_mem(param_18, param_19, param_20, v_260, v_260BufferSize); + write_mem(param_18, param_19, param_20, v_266, v_266BufferSize); Alloc param_21 = a; uint param_22 = ix + 7u; uint param_23 = as_type(s.c1.x); - write_mem(param_21, param_22, param_23, v_260, v_260BufferSize); + write_mem(param_21, param_22, param_23, v_266, v_266BufferSize); Alloc param_24 = a; uint param_25 = ix + 8u; uint param_26 = as_type(s.c1.y); - write_mem(param_24, param_25, param_26, v_260, v_260BufferSize); + write_mem(param_24, param_25, param_26, v_266, v_266BufferSize); Alloc param_27 = a; uint param_28 = ix + 9u; uint param_29 = as_type(s.ra); - write_mem(param_27, param_28, param_29, v_260, v_260BufferSize); + write_mem(param_27, param_28, param_29, v_266, v_266BufferSize); Alloc param_30 = a; uint param_31 = ix + 10u; uint param_32 = as_type(s.roff); - write_mem(param_30, param_31, param_32, v_260, v_260BufferSize); + write_mem(param_30, param_31, param_32, v_266, v_266BufferSize); } static inline __attribute__((always_inline)) -void Cmd_RadGrad_write(thread const Alloc& a, thread const CmdRef& ref, thread const CmdRadGrad& s, device Memory& v_260, constant uint& v_260BufferSize) +void Cmd_RadGrad_write(thread const Alloc& a, thread const CmdRef& ref, thread const CmdRadGrad& s, device Memory& v_266, constant uint& v_266BufferSize) { Alloc param = a; uint param_1 = ref.offset >> uint(2); uint param_2 = 7u; - write_mem(param, param_1, param_2, v_260, v_260BufferSize); + write_mem(param, param_1, param_2, v_266, v_266BufferSize); Alloc param_3 = a; CmdRadGradRef param_4 = CmdRadGradRef{ ref.offset + 4u }; CmdRadGrad param_5 = s; - CmdRadGrad_write(param_3, param_4, param_5, v_260, v_260BufferSize); + CmdRadGrad_write(param_3, param_4, param_5, v_266, v_266BufferSize); } static inline __attribute__((always_inline)) -void CmdImage_write(thread const Alloc& a, thread const CmdImageRef& ref, thread const CmdImage& s, device Memory& v_260, constant uint& v_260BufferSize) +void CmdImage_write(thread const Alloc& a, thread const CmdImageRef& ref, thread const CmdImage& s, device Memory& v_266, constant uint& v_266BufferSize) { uint ix = ref.offset >> uint(2); Alloc param = a; uint param_1 = ix + 0u; uint param_2 = s.index; - write_mem(param, param_1, param_2, v_260, v_260BufferSize); + write_mem(param, param_1, param_2, v_266, v_266BufferSize); Alloc param_3 = a; uint param_4 = ix + 1u; uint param_5 = (uint(s.offset.x) & 65535u) | (uint(s.offset.y) << uint(16)); - write_mem(param_3, param_4, param_5, v_260, v_260BufferSize); + write_mem(param_3, param_4, param_5, v_266, v_266BufferSize); } static inline __attribute__((always_inline)) -void Cmd_Image_write(thread const Alloc& a, thread const CmdRef& ref, thread const CmdImage& s, device Memory& v_260, constant uint& v_260BufferSize) +void Cmd_Image_write(thread const Alloc& a, thread const CmdRef& ref, thread const CmdImage& s, device Memory& v_266, constant uint& v_266BufferSize) { Alloc param = a; uint param_1 = ref.offset >> uint(2); uint param_2 = 8u; - write_mem(param, param_1, param_2, v_260, v_260BufferSize); + write_mem(param, param_1, param_2, v_266, v_266BufferSize); Alloc param_3 = a; CmdImageRef param_4 = CmdImageRef{ ref.offset + 4u }; CmdImage param_5 = s; - CmdImage_write(param_3, param_4, param_5, v_260, v_260BufferSize); + CmdImage_write(param_3, param_4, param_5, v_266, v_266BufferSize); } static inline __attribute__((always_inline)) -void Cmd_BeginClip_write(thread const Alloc& a, thread const CmdRef& ref, device Memory& v_260, constant uint& v_260BufferSize) +void Cmd_BeginClip_write(thread const Alloc& a, thread const CmdRef& ref, device Memory& v_266, constant uint& v_266BufferSize) { Alloc param = a; uint param_1 = ref.offset >> uint(2); uint param_2 = 9u; - write_mem(param, param_1, param_2, v_260, v_260BufferSize); + write_mem(param, param_1, param_2, v_266, v_266BufferSize); } static inline __attribute__((always_inline)) -void CmdEndClip_write(thread const Alloc& a, thread const CmdEndClipRef& ref, thread const CmdEndClip& s, device Memory& v_260, constant uint& v_260BufferSize) +void CmdEndClip_write(thread const Alloc& a, thread const CmdEndClipRef& ref, thread const CmdEndClip& s, device Memory& v_266, constant uint& v_266BufferSize) { 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, v_260, v_260BufferSize); + write_mem(param, param_1, param_2, v_266, v_266BufferSize); } static inline __attribute__((always_inline)) -void Cmd_EndClip_write(thread const Alloc& a, thread const CmdRef& ref, thread const CmdEndClip& s, device Memory& v_260, constant uint& v_260BufferSize) +void Cmd_EndClip_write(thread const Alloc& a, thread const CmdRef& ref, thread const CmdEndClip& s, device Memory& v_266, constant uint& v_266BufferSize) { Alloc param = a; uint param_1 = ref.offset >> uint(2); uint param_2 = 10u; - write_mem(param, param_1, param_2, v_260, v_260BufferSize); + write_mem(param, param_1, param_2, v_266, v_266BufferSize); Alloc param_3 = a; CmdEndClipRef param_4 = CmdEndClipRef{ ref.offset + 4u }; CmdEndClip param_5 = s; - CmdEndClip_write(param_3, param_4, param_5, v_260, v_260BufferSize); + CmdEndClip_write(param_3, param_4, param_5, v_266, v_266BufferSize); } static inline __attribute__((always_inline)) -void Cmd_End_write(thread const Alloc& a, thread const CmdRef& ref, device Memory& v_260, constant uint& v_260BufferSize) +void Cmd_End_write(thread const Alloc& a, thread const CmdRef& ref, device Memory& v_266, constant uint& v_266BufferSize) { Alloc param = a; uint param_1 = ref.offset >> uint(2); uint param_2 = 0u; - write_mem(param, param_1, param_2, v_260, v_260BufferSize); + write_mem(param, param_1, param_2, v_266, v_266BufferSize); } -kernel void main0(constant uint* spvBufferSizeConstants [[buffer(25)]], device Memory& v_260 [[buffer(0)]], const device ConfigBuf& _1005 [[buffer(1)]], const device SceneBuf& _1378 [[buffer(2)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]]) +static inline __attribute__((always_inline)) +void alloc_write(thread const Alloc& a, thread const uint& offset, thread const Alloc& alloc, device Memory& v_266, constant uint& v_266BufferSize) +{ + Alloc param = a; + uint param_1 = offset >> uint(2); + uint param_2 = alloc.offset; + write_mem(param, param_1, param_2, v_266, v_266BufferSize); +} + +kernel void main0(constant uint* spvBufferSizeConstants [[buffer(25)]], device Memory& v_266 [[buffer(0)]], const device ConfigBuf& _1020 [[buffer(1)]], const device SceneBuf& _1399 [[buffer(2)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]]) { threadgroup uint sh_bitmaps[8][256]; threadgroup Alloc sh_part_elements[256]; @@ -689,19 +698,19 @@ kernel void main0(constant uint* spvBufferSizeConstants [[buffer(25)]], device M threadgroup uint sh_tile_y0[256]; threadgroup uint sh_tile_base[256]; threadgroup uint sh_tile_count[256]; - constant uint& v_260BufferSize = spvBufferSizeConstants[0]; - uint width_in_bins = ((_1005.conf.width_in_tiles + 16u) - 1u) / 16u; + constant uint& v_266BufferSize = spvBufferSizeConstants[0]; + uint width_in_bins = ((_1020.conf.width_in_tiles + 16u) - 1u) / 16u; uint bin_ix = (width_in_bins * gl_WorkGroupID.y) + gl_WorkGroupID.x; uint partition_ix = 0u; - uint n_partitions = ((_1005.conf.n_elements + 256u) - 1u) / 256u; + uint n_partitions = ((_1020.conf.n_elements + 256u) - 1u) / 256u; uint th_ix = gl_LocalInvocationID.x; uint bin_tile_x = 16u * gl_WorkGroupID.x; uint bin_tile_y = 16u * gl_WorkGroupID.y; uint tile_x = gl_LocalInvocationID.x % 16u; uint tile_y = gl_LocalInvocationID.x / 16u; - uint this_tile_ix = (((bin_tile_y + tile_y) * _1005.conf.width_in_tiles) + bin_tile_x) + tile_x; + uint this_tile_ix = (((bin_tile_y + tile_y) * _1020.conf.width_in_tiles) + bin_tile_x) + tile_x; Alloc param; - param.offset = _1005.conf.ptcl_alloc.offset; + param.offset = _1020.conf.ptcl_alloc.offset; uint param_1 = this_tile_ix * 1024u; uint param_2 = 1024u; Alloc cmd_alloc = slice_mem(param, param_1, param_2); @@ -713,21 +722,25 @@ kernel void main0(constant uint* spvBufferSizeConstants [[buffer(25)]], device M uint wr_ix = 0u; uint part_start_ix = 0u; uint ready_ix = 0u; + Alloc param_3 = cmd_alloc; + uint param_4 = 0u; + uint param_5 = 8u; + Alloc scratch_alloc = slice_mem(param_3, param_4, param_5); cmd_ref.offset += 4u; uint render_blend_depth = 0u; uint max_blend_depth = 0u; - uint drawmonoid_start = _1005.conf.drawmonoid_alloc.offset >> uint(2); - uint drawtag_start = _1005.conf.drawtag_offset >> uint(2); - uint drawdata_start = _1005.conf.drawdata_offset >> uint(2); - uint drawinfo_start = _1005.conf.drawinfo_alloc.offset >> uint(2); - bool mem_ok = v_260.mem_error == 0u; - Alloc param_3; - Alloc param_5; - uint _1310; + uint drawmonoid_start = _1020.conf.drawmonoid_alloc.offset >> uint(2); + uint drawtag_start = _1020.conf.drawtag_offset >> uint(2); + uint drawdata_start = _1020.conf.drawdata_offset >> uint(2); + uint drawinfo_start = _1020.conf.drawinfo_alloc.offset >> uint(2); + bool mem_ok = v_266.mem_error == 0u; + Alloc param_6; + Alloc param_8; + uint _1331; uint element_ix; - Alloc param_14; + Alloc param_17; uint tile_count; - uint _1611; + uint _1632; float linewidth; CmdLinGrad cmd_lin; CmdRadGrad cmd_rad; @@ -737,36 +750,36 @@ kernel void main0(constant uint* spvBufferSizeConstants [[buffer(25)]], device M { sh_bitmaps[i][th_ix] = 0u; } - bool _1362; + bool _1383; for (;;) { if ((ready_ix == wr_ix) && (partition_ix < n_partitions)) { part_start_ix = ready_ix; uint count = 0u; - bool _1160 = th_ix < 256u; - bool _1168; - if (_1160) + bool _1181 = th_ix < 256u; + bool _1189; + if (_1181) { - _1168 = (partition_ix + th_ix) < n_partitions; + _1189 = (partition_ix + th_ix) < n_partitions; } else { - _1168 = _1160; + _1189 = _1181; } - if (_1168) + if (_1189) { - uint in_ix = (_1005.conf.bin_alloc.offset >> uint(2)) + ((((partition_ix + th_ix) * 256u) + bin_ix) * 2u); - param_3.offset = _1005.conf.bin_alloc.offset; - uint param_4 = in_ix; - count = read_mem(param_3, param_4, v_260, v_260BufferSize); - param_5.offset = _1005.conf.bin_alloc.offset; - uint param_6 = in_ix + 1u; - uint offset = read_mem(param_5, param_6, v_260, v_260BufferSize); - uint param_7 = offset; - uint param_8 = count * 4u; - bool param_9 = mem_ok; - sh_part_elements[th_ix] = new_alloc(param_7, param_8, param_9); + uint in_ix = (_1020.conf.bin_alloc.offset >> uint(2)) + ((((partition_ix + th_ix) * 256u) + bin_ix) * 2u); + param_6.offset = _1020.conf.bin_alloc.offset; + uint param_7 = in_ix; + count = read_mem(param_6, param_7, v_266, v_266BufferSize); + param_8.offset = _1020.conf.bin_alloc.offset; + uint param_9 = in_ix + 1u; + uint offset = read_mem(param_8, param_9, v_266, v_266BufferSize); + uint param_10 = offset; + uint param_11 = count * 4u; + bool param_12 = mem_ok; + sh_part_elements[th_ix] = new_alloc(param_10, param_11, param_12); } for (uint i_1 = 0u; i_1 < 8u; i_1++) { @@ -806,34 +819,34 @@ kernel void main0(constant uint* spvBufferSizeConstants [[buffer(25)]], device M } if (part_ix > 0u) { - _1310 = sh_part_count[part_ix - 1u]; + _1331 = sh_part_count[part_ix - 1u]; } else { - _1310 = part_start_ix; + _1331 = part_start_ix; } - ix -= _1310; + ix -= _1331; Alloc bin_alloc = sh_part_elements[part_ix]; BinInstanceRef inst_ref = BinInstanceRef{ bin_alloc.offset }; - BinInstanceRef param_10 = inst_ref; - uint param_11 = ix; - Alloc param_12 = bin_alloc; - BinInstanceRef param_13 = BinInstance_index(param_10, param_11); - BinInstance inst = BinInstance_read(param_12, param_13, v_260, v_260BufferSize); + BinInstanceRef param_13 = inst_ref; + uint param_14 = ix; + Alloc param_15 = bin_alloc; + BinInstanceRef param_16 = BinInstance_index(param_13, param_14); + BinInstance inst = BinInstance_read(param_15, param_16, v_266, v_266BufferSize); sh_elements[th_ix] = inst.element_ix; } threadgroup_barrier(mem_flags::mem_threadgroup); wr_ix = min((rd_ix + 256u), ready_ix); - bool _1352 = (wr_ix - rd_ix) < 256u; - if (_1352) + bool _1373 = (wr_ix - rd_ix) < 256u; + if (_1373) { - _1362 = (wr_ix < ready_ix) || (partition_ix < n_partitions); + _1383 = (wr_ix < ready_ix) || (partition_ix < n_partitions); } else { - _1362 = _1352; + _1383 = _1373; } - if (_1362) + if (_1383) { continue; } @@ -846,7 +859,7 @@ kernel void main0(constant uint* spvBufferSizeConstants [[buffer(25)]], device M if ((th_ix + rd_ix) < wr_ix) { element_ix = sh_elements[th_ix]; - tag = _1378.scene[drawtag_start + element_ix]; + tag = _1399.scene[drawtag_start + element_ix]; } switch (tag) { @@ -858,10 +871,10 @@ kernel void main0(constant uint* spvBufferSizeConstants [[buffer(25)]], device M case 37u: { uint drawmonoid_base = drawmonoid_start + (4u * element_ix); - uint path_ix = v_260.memory[drawmonoid_base]; - param_14.offset = _1005.conf.tile_alloc.offset; - PathRef param_15 = PathRef{ _1005.conf.tile_alloc.offset + (path_ix * 12u) }; - Path path = Path_read(param_14, param_15, v_260, v_260BufferSize); + uint path_ix = v_266.memory[drawmonoid_base]; + param_17.offset = _1020.conf.tile_alloc.offset; + PathRef param_18 = PathRef{ _1020.conf.tile_alloc.offset + (path_ix * 12u) }; + Path path = Path_read(param_17, param_18, v_266, v_266BufferSize); uint stride = path.bbox.z - path.bbox.x; sh_tile_stride[th_ix] = stride; int dx = int(path.bbox.x) - int(bin_tile_x); @@ -876,13 +889,13 @@ kernel void main0(constant uint* spvBufferSizeConstants [[buffer(25)]], device M tile_count = uint(x1 - x0) * uint(y1 - y0); uint base = path.tiles.offset - (((uint(dy) * stride) + uint(dx)) * 8u); sh_tile_base[th_ix] = base; - uint param_16 = path.tiles.offset; - uint param_17 = ((path.bbox.z - path.bbox.x) * (path.bbox.w - path.bbox.y)) * 8u; - bool param_18 = mem_ok; - Alloc path_alloc = new_alloc(param_16, param_17, param_18); - uint param_19 = th_ix; - Alloc param_20 = path_alloc; - write_tile_alloc(param_19, param_20); + uint param_19 = path.tiles.offset; + uint param_20 = ((path.bbox.z - path.bbox.x) * (path.bbox.w - path.bbox.y)) * 8u; + bool param_21 = mem_ok; + Alloc path_alloc = new_alloc(param_19, param_20, param_21); + uint param_22 = th_ix; + Alloc param_23 = path_alloc; + write_tile_alloc(param_22, param_23); break; } default: @@ -916,54 +929,54 @@ kernel void main0(constant uint* spvBufferSizeConstants [[buffer(25)]], device M } } uint element_ix_1 = sh_elements[el_ix]; - uint tag_1 = _1378.scene[drawtag_start + element_ix_1]; + uint tag_1 = _1399.scene[drawtag_start + element_ix_1]; if (el_ix > 0u) { - _1611 = sh_tile_count[el_ix - 1u]; + _1632 = sh_tile_count[el_ix - 1u]; } else { - _1611 = 0u; + _1632 = 0u; } - uint seq_ix = ix_1 - _1611; + uint seq_ix = ix_1 - _1632; uint width = sh_tile_width[el_ix]; uint x = sh_tile_x0[el_ix] + (seq_ix % width); uint y = sh_tile_y0[el_ix] + (seq_ix / width); bool include_tile = false; if (mem_ok) { - uint param_21 = el_ix; - bool param_22 = mem_ok; - Alloc param_23 = read_tile_alloc(param_21, param_22, v_260, v_260BufferSize); - TileRef param_24 = TileRef{ sh_tile_base[el_ix] + (((sh_tile_stride[el_ix] * y) + x) * 8u) }; - Tile tile = Tile_read(param_23, param_24, v_260, v_260BufferSize); + uint param_24 = el_ix; + bool param_25 = mem_ok; + Alloc param_26 = read_tile_alloc(param_24, param_25, v_266, v_266BufferSize); + TileRef param_27 = TileRef{ sh_tile_base[el_ix] + (((sh_tile_stride[el_ix] * y) + x) * 8u) }; + Tile tile = Tile_read(param_26, param_27, v_266, v_266BufferSize); bool is_clip = (tag_1 & 1u) != 0u; bool is_blend = false; if (is_clip) { uint drawmonoid_base_1 = drawmonoid_start + (4u * element_ix_1); - uint scene_offset = v_260.memory[drawmonoid_base_1 + 2u]; + uint scene_offset = v_266.memory[drawmonoid_base_1 + 2u]; uint dd = drawdata_start + (scene_offset >> uint(2)); - uint blend = _1378.scene[dd]; + uint blend = _1399.scene[dd]; is_blend = blend != 32771u; } - bool _1699 = tile.tile.offset != 0u; - bool _1708; - if (!_1699) + bool _1720 = tile.tile.offset != 0u; + bool _1729; + if (!_1720) { - _1708 = (tile.backdrop == 0) == is_clip; + _1729 = (tile.backdrop == 0) == is_clip; } else { - _1708 = _1699; + _1729 = _1720; } - include_tile = _1708 || is_blend; + include_tile = _1729 || is_blend; } if (include_tile) { uint el_slice = el_ix / 32u; uint el_mask = 1u << (el_ix & 31u); - uint _1730 = atomic_fetch_or_explicit((threadgroup atomic_uint*)&sh_bitmaps[el_slice][(y * 16u) + x], el_mask, memory_order_relaxed); + uint _1751 = atomic_fetch_or_explicit((threadgroup atomic_uint*)&sh_bitmaps[el_slice][(y * 16u) + x], el_mask, memory_order_relaxed); } } threadgroup_barrier(mem_flags::mem_threadgroup); @@ -987,175 +1000,175 @@ kernel void main0(constant uint* spvBufferSizeConstants [[buffer(25)]], device M uint element_ref_ix = (slice_ix * 32u) + uint(int(spvFindLSB(bitmap))); uint element_ix_2 = sh_elements[element_ref_ix]; bitmap &= (bitmap - 1u); - uint drawtag = _1378.scene[drawtag_start + element_ix_2]; + uint drawtag = _1399.scene[drawtag_start + element_ix_2]; if (clip_zero_depth == 0u) { - uint param_25 = element_ref_ix; - bool param_26 = mem_ok; - Alloc param_27 = read_tile_alloc(param_25, param_26, v_260, v_260BufferSize); - TileRef param_28 = TileRef{ sh_tile_base[element_ref_ix] + (((sh_tile_stride[element_ref_ix] * tile_y) + tile_x) * 8u) }; - Tile tile_1 = Tile_read(param_27, param_28, v_260, v_260BufferSize); + uint param_28 = element_ref_ix; + bool param_29 = mem_ok; + Alloc param_30 = read_tile_alloc(param_28, param_29, v_266, v_266BufferSize); + TileRef param_31 = TileRef{ sh_tile_base[element_ref_ix] + (((sh_tile_stride[element_ref_ix] * tile_y) + tile_x) * 8u) }; + Tile tile_1 = Tile_read(param_30, param_31, v_266, v_266BufferSize); uint drawmonoid_base_2 = drawmonoid_start + (4u * element_ix_2); - uint scene_offset_1 = v_260.memory[drawmonoid_base_2 + 2u]; - uint info_offset = v_260.memory[drawmonoid_base_2 + 3u]; + uint scene_offset_1 = v_266.memory[drawmonoid_base_2 + 2u]; + uint info_offset = v_266.memory[drawmonoid_base_2 + 3u]; uint dd_1 = drawdata_start + (scene_offset_1 >> uint(2)); uint di = drawinfo_start + (info_offset >> uint(2)); switch (drawtag) { case 68u: { - linewidth = as_type(v_260.memory[di]); - Alloc param_29 = cmd_alloc; - CmdRef param_30 = cmd_ref; - uint param_31 = cmd_limit; - bool _1855 = alloc_cmd(param_29, param_30, param_31, v_260, v_260BufferSize); - cmd_alloc = param_29; - cmd_ref = param_30; - cmd_limit = param_31; - if (!_1855) + linewidth = as_type(v_266.memory[di]); + Alloc param_32 = cmd_alloc; + CmdRef param_33 = cmd_ref; + uint param_34 = cmd_limit; + bool _1876 = alloc_cmd(param_32, param_33, param_34, v_266, v_266BufferSize); + cmd_alloc = param_32; + cmd_ref = param_33; + cmd_limit = param_34; + if (!_1876) { break; } - Alloc param_32 = cmd_alloc; - CmdRef param_33 = cmd_ref; - Tile param_34 = tile_1; - float param_35 = linewidth; - write_fill(param_32, param_33, param_34, param_35, v_260, v_260BufferSize); - cmd_ref = param_33; - uint rgba = _1378.scene[dd_1]; - Alloc param_36 = cmd_alloc; - CmdRef param_37 = cmd_ref; - CmdColor param_38 = CmdColor{ rgba }; - Cmd_Color_write(param_36, param_37, param_38, v_260, v_260BufferSize); + Alloc param_35 = cmd_alloc; + CmdRef param_36 = cmd_ref; + Tile param_37 = tile_1; + float param_38 = linewidth; + write_fill(param_35, param_36, param_37, param_38, v_266, v_266BufferSize); + cmd_ref = param_36; + uint rgba = _1399.scene[dd_1]; + Alloc param_39 = cmd_alloc; + CmdRef param_40 = cmd_ref; + CmdColor param_41 = CmdColor{ rgba }; + Cmd_Color_write(param_39, param_40, param_41, v_266, v_266BufferSize); cmd_ref.offset += 8u; break; } case 276u: { - Alloc param_39 = cmd_alloc; - CmdRef param_40 = cmd_ref; - uint param_41 = cmd_limit; - bool _1896 = alloc_cmd(param_39, param_40, param_41, v_260, v_260BufferSize); - cmd_alloc = param_39; - cmd_ref = param_40; - cmd_limit = param_41; - if (!_1896) + Alloc param_42 = cmd_alloc; + CmdRef param_43 = cmd_ref; + uint param_44 = cmd_limit; + bool _1917 = alloc_cmd(param_42, param_43, param_44, v_266, v_266BufferSize); + cmd_alloc = param_42; + cmd_ref = param_43; + cmd_limit = param_44; + if (!_1917) { break; } - linewidth = as_type(v_260.memory[di]); - Alloc param_42 = cmd_alloc; - CmdRef param_43 = cmd_ref; - Tile param_44 = tile_1; - float param_45 = linewidth; - write_fill(param_42, param_43, param_44, param_45, v_260, v_260BufferSize); - cmd_ref = param_43; - cmd_lin.index = _1378.scene[dd_1]; - cmd_lin.line_x = as_type(v_260.memory[di + 1u]); - cmd_lin.line_y = as_type(v_260.memory[di + 2u]); - cmd_lin.line_c = as_type(v_260.memory[di + 3u]); - Alloc param_46 = cmd_alloc; - CmdRef param_47 = cmd_ref; - CmdLinGrad param_48 = cmd_lin; - Cmd_LinGrad_write(param_46, param_47, param_48, v_260, v_260BufferSize); + linewidth = as_type(v_266.memory[di]); + Alloc param_45 = cmd_alloc; + CmdRef param_46 = cmd_ref; + Tile param_47 = tile_1; + float param_48 = linewidth; + write_fill(param_45, param_46, param_47, param_48, v_266, v_266BufferSize); + cmd_ref = param_46; + cmd_lin.index = _1399.scene[dd_1]; + cmd_lin.line_x = as_type(v_266.memory[di + 1u]); + cmd_lin.line_y = as_type(v_266.memory[di + 2u]); + cmd_lin.line_c = as_type(v_266.memory[di + 3u]); + Alloc param_49 = cmd_alloc; + CmdRef param_50 = cmd_ref; + CmdLinGrad param_51 = cmd_lin; + Cmd_LinGrad_write(param_49, param_50, param_51, v_266, v_266BufferSize); cmd_ref.offset += 20u; break; } case 732u: { - Alloc param_49 = cmd_alloc; - CmdRef param_50 = cmd_ref; - uint param_51 = cmd_limit; - bool _1960 = alloc_cmd(param_49, param_50, param_51, v_260, v_260BufferSize); - cmd_alloc = param_49; - cmd_ref = param_50; - cmd_limit = param_51; - if (!_1960) + Alloc param_52 = cmd_alloc; + CmdRef param_53 = cmd_ref; + uint param_54 = cmd_limit; + bool _1981 = alloc_cmd(param_52, param_53, param_54, v_266, v_266BufferSize); + cmd_alloc = param_52; + cmd_ref = param_53; + cmd_limit = param_54; + if (!_1981) { break; } - linewidth = as_type(v_260.memory[di]); - Alloc param_52 = cmd_alloc; - CmdRef param_53 = cmd_ref; - Tile param_54 = tile_1; - float param_55 = linewidth; - write_fill(param_52, param_53, param_54, param_55, v_260, v_260BufferSize); - cmd_ref = param_53; - cmd_rad.index = _1378.scene[dd_1]; - cmd_rad.mat = as_type(uint4(v_260.memory[di + 1u], v_260.memory[di + 2u], v_260.memory[di + 3u], v_260.memory[di + 4u])); - cmd_rad.xlat = as_type(uint2(v_260.memory[di + 5u], v_260.memory[di + 6u])); - cmd_rad.c1 = as_type(uint2(v_260.memory[di + 7u], v_260.memory[di + 8u])); - cmd_rad.ra = as_type(v_260.memory[di + 9u]); - cmd_rad.roff = as_type(v_260.memory[di + 10u]); - Alloc param_56 = cmd_alloc; - CmdRef param_57 = cmd_ref; - CmdRadGrad param_58 = cmd_rad; - Cmd_RadGrad_write(param_56, param_57, param_58, v_260, v_260BufferSize); + linewidth = as_type(v_266.memory[di]); + Alloc param_55 = cmd_alloc; + CmdRef param_56 = cmd_ref; + Tile param_57 = tile_1; + float param_58 = linewidth; + write_fill(param_55, param_56, param_57, param_58, v_266, v_266BufferSize); + cmd_ref = param_56; + cmd_rad.index = _1399.scene[dd_1]; + cmd_rad.mat = as_type(uint4(v_266.memory[di + 1u], v_266.memory[di + 2u], v_266.memory[di + 3u], v_266.memory[di + 4u])); + cmd_rad.xlat = as_type(uint2(v_266.memory[di + 5u], v_266.memory[di + 6u])); + cmd_rad.c1 = as_type(uint2(v_266.memory[di + 7u], v_266.memory[di + 8u])); + cmd_rad.ra = as_type(v_266.memory[di + 9u]); + cmd_rad.roff = as_type(v_266.memory[di + 10u]); + Alloc param_59 = cmd_alloc; + CmdRef param_60 = cmd_ref; + CmdRadGrad param_61 = cmd_rad; + Cmd_RadGrad_write(param_59, param_60, param_61, v_266, v_266BufferSize); cmd_ref.offset += 48u; break; } case 72u: { - linewidth = as_type(v_260.memory[di]); - Alloc param_59 = cmd_alloc; - CmdRef param_60 = cmd_ref; - uint param_61 = cmd_limit; - bool _2066 = alloc_cmd(param_59, param_60, param_61, v_260, v_260BufferSize); - cmd_alloc = param_59; - cmd_ref = param_60; - cmd_limit = param_61; - if (!_2066) + linewidth = as_type(v_266.memory[di]); + Alloc param_62 = cmd_alloc; + CmdRef param_63 = cmd_ref; + uint param_64 = cmd_limit; + bool _2087 = alloc_cmd(param_62, param_63, param_64, v_266, v_266BufferSize); + cmd_alloc = param_62; + cmd_ref = param_63; + cmd_limit = param_64; + if (!_2087) { break; } - Alloc param_62 = cmd_alloc; - CmdRef param_63 = cmd_ref; - Tile param_64 = tile_1; - float param_65 = linewidth; - write_fill(param_62, param_63, param_64, param_65, v_260, v_260BufferSize); - cmd_ref = param_63; - uint index = _1378.scene[dd_1]; - uint raw1 = _1378.scene[dd_1 + 1u]; + Alloc param_65 = cmd_alloc; + CmdRef param_66 = cmd_ref; + Tile param_67 = tile_1; + float param_68 = linewidth; + write_fill(param_65, param_66, param_67, param_68, v_266, v_266BufferSize); + cmd_ref = param_66; + uint index = _1399.scene[dd_1]; + uint raw1 = _1399.scene[dd_1 + 1u]; int2 offset_1 = int2(int(raw1 << uint(16)) >> 16, int(raw1) >> 16); - Alloc param_66 = cmd_alloc; - CmdRef param_67 = cmd_ref; - CmdImage param_68 = CmdImage{ index, offset_1 }; - Cmd_Image_write(param_66, param_67, param_68, v_260, v_260BufferSize); + Alloc param_69 = cmd_alloc; + CmdRef param_70 = cmd_ref; + CmdImage param_71 = CmdImage{ index, offset_1 }; + Cmd_Image_write(param_69, param_70, param_71, v_266, v_266BufferSize); cmd_ref.offset += 12u; break; } case 5u: { - bool _2119 = tile_1.tile.offset == 0u; - bool _2125; - if (_2119) + bool _2140 = tile_1.tile.offset == 0u; + bool _2146; + if (_2140) { - _2125 = tile_1.backdrop == 0; + _2146 = tile_1.backdrop == 0; } else { - _2125 = _2119; + _2146 = _2140; } - if (_2125) + if (_2146) { clip_zero_depth = clip_depth + 1u; } else { - Alloc param_69 = cmd_alloc; - CmdRef param_70 = cmd_ref; - uint param_71 = cmd_limit; - bool _2137 = alloc_cmd(param_69, param_70, param_71, v_260, v_260BufferSize); - cmd_alloc = param_69; - cmd_ref = param_70; - cmd_limit = param_71; - if (!_2137) + Alloc param_72 = cmd_alloc; + CmdRef param_73 = cmd_ref; + uint param_74 = cmd_limit; + bool _2158 = alloc_cmd(param_72, param_73, param_74, v_266, v_266BufferSize); + cmd_alloc = param_72; + cmd_ref = param_73; + cmd_limit = param_74; + if (!_2158) { break; } - Alloc param_72 = cmd_alloc; - CmdRef param_73 = cmd_ref; - Cmd_BeginClip_write(param_72, param_73, v_260, v_260BufferSize); + Alloc param_75 = cmd_alloc; + CmdRef param_76 = cmd_ref; + Cmd_BeginClip_write(param_75, param_76, v_266, v_266BufferSize); cmd_ref.offset += 4u; render_blend_depth++; max_blend_depth = max(max_blend_depth, render_blend_depth); @@ -1166,28 +1179,28 @@ kernel void main0(constant uint* spvBufferSizeConstants [[buffer(25)]], device M case 37u: { clip_depth--; - Alloc param_74 = cmd_alloc; - CmdRef param_75 = cmd_ref; - uint param_76 = cmd_limit; - bool _2170 = alloc_cmd(param_74, param_75, param_76, v_260, v_260BufferSize); - cmd_alloc = param_74; - cmd_ref = param_75; - cmd_limit = param_76; - if (!_2170) + Alloc param_77 = cmd_alloc; + CmdRef param_78 = cmd_ref; + uint param_79 = cmd_limit; + bool _2191 = alloc_cmd(param_77, param_78, param_79, v_266, v_266BufferSize); + cmd_alloc = param_77; + cmd_ref = param_78; + cmd_limit = param_79; + if (!_2191) { break; } - Alloc param_77 = cmd_alloc; - CmdRef param_78 = cmd_ref; - Tile param_79 = tile_1; - float param_80 = -1.0; - write_fill(param_77, param_78, param_79, param_80, v_260, v_260BufferSize); - cmd_ref = param_78; - uint blend_1 = _1378.scene[dd_1]; - Alloc param_81 = cmd_alloc; - CmdRef param_82 = cmd_ref; - CmdEndClip param_83 = CmdEndClip{ blend_1 }; - Cmd_EndClip_write(param_81, param_82, param_83, v_260, v_260BufferSize); + Alloc param_80 = cmd_alloc; + CmdRef param_81 = cmd_ref; + Tile param_82 = tile_1; + float param_83 = -1.0; + write_fill(param_80, param_81, param_82, param_83, v_266, v_266BufferSize); + cmd_ref = param_81; + uint blend_1 = _1399.scene[dd_1]; + Alloc param_84 = cmd_alloc; + CmdRef param_85 = cmd_ref; + CmdEndClip param_86 = CmdEndClip{ blend_1 }; + Cmd_EndClip_write(param_84, param_85, param_86, v_266, v_266BufferSize); cmd_ref.offset += 8u; render_blend_depth--; break; @@ -1222,23 +1235,31 @@ kernel void main0(constant uint* spvBufferSizeConstants [[buffer(25)]], device M break; } } - bool _2242 = (bin_tile_x + tile_x) < _1005.conf.width_in_tiles; - bool _2251; - if (_2242) + bool _2263 = (bin_tile_x + tile_x) < _1020.conf.width_in_tiles; + bool _2272; + if (_2263) { - _2251 = (bin_tile_y + tile_y) < _1005.conf.height_in_tiles; + _2272 = (bin_tile_y + tile_y) < _1020.conf.height_in_tiles; } else { - _2251 = _2242; + _2272 = _2263; } - if (_2251) + if (_2272) { - Alloc param_84 = cmd_alloc; - CmdRef param_85 = cmd_ref; - Cmd_End_write(param_84, param_85, v_260, v_260BufferSize); + Alloc param_87 = cmd_alloc; + CmdRef param_88 = cmd_ref; + Cmd_End_write(param_87, param_88, v_266, v_266BufferSize); if (max_blend_depth > 4u) { + uint scratch_size = (((max_blend_depth * 16u) * 16u) * 1u) * 4u; + uint param_89 = scratch_size; + MallocResult _2293 = malloc(param_89, v_266, v_266BufferSize); + MallocResult scratch = _2293; + Alloc param_90 = scratch_alloc; + uint param_91 = scratch_alloc.offset; + Alloc param_92 = scratch.alloc; + alloc_write(param_90, param_91, param_92, v_266, v_266BufferSize); } } } diff --git a/piet-gpu/shader/gen/coarse.spv b/piet-gpu/shader/gen/coarse.spv index bcac8449768a7c22e0e02b9d59f6cdd0f89ee6e2..fe5eeee188be90b85271b37aad93d5f1f23c237f 100644 GIT binary patch literal 60516 zcmbWg1)wFx)wSJl?+oq^2@LK83lbp60E4?E*l-zPE{rjP2X}Y(0E1g_5+HbjTOe4F z5S&0DXbAuFoYQM=-A?l5{d3;8Q){igYuBz_)qVP$K6mJtbpFYzYLaToYVxX@qH4Wn zs3t|Js;R2Jt#-ghTW++}*l~lGTK>mN>o9HA(el%0x@xkj4{ayMh_3FD8YY9x&M`gx z)n`z|S$~fX{%=y+`yic+JYb`N12)}sz((r~95j0Lz_H`@89rp-@a~aA1`h5XGi1=X z?mdU-H+kXLJ#x_S@qjT&?^DeTA2BhJ_82pv^^7@XH79)BsPTh_4P@j+#*P~^ ze$coj*EIPh$1l^$zTIzj{4$=o&H?Y4_f%Cs_+DeW#|`<9aVP{n_WdDk8EbBMPppn= z9{8SJ!||Q6S`a*DNY~*17#6!R=cg@WE&%U|IaRedeC+Vdt-@2&5Q)lYl-4s`u8{F)QDR$Y~UlaDdYXI0vEeLAYu zz{zb5@Mc}(hFJv9L%FY>8gy3u;ahPk#7FL{^jRBj9ChQ!eeJT#C}z^5O#&WAtg*Cx zI7`=~wcr1HVs(bC3b*I8F8wl}0pRw0`cxZ&&2`HmyP9q*7d2mHz8k~!Q#W6^FLSeX zP1MYHm##s(4<0jWw7NNOPHW%)cg}i{LXq>P^h?g0f!jIvt+oXBA2E2#abrg9-rkMv zoYZ~w0GO~)I`5b zrKvN&C5Li#Bc}W8=d_vQcHs6LJF4x$?KSUH?F1(7y4}Nv6SX~8HD5h*?W}f&*M6;e z%6&EGxii|ouG1`)g{b@g_gU82w3e%9j;hYE+Zj5)LB!1bcx#wAKQ~S0=Sjf0{YMQS zHO2|Lzgi~-de8j&RJ#>^z}6hqe2v+4m#%?>j9^`7soZ>t`QIm#p&PSl$t|rlml4Fw zT)a8RT8#p?=hC-gc*otSd*pyIU4vW8)tZmGuW~-VsrLA`*R}OVi=OMBwNP`tTK5rL z9in;P(bD$m(K@10_&&T-?hEf3G=BuY}ZOyv%^qsXj8J@oXea1R+ntd*9kKccvQP1FwIn{=i+KSy#okg2Doej>M&H-o5&IdF82IEJxC*)eF`RbWZ zXLTXG_G_(|+?V(9?YOg4m%(eF|2~m=#KtUCGX*l~s4gL9=I~2!=5Q&vJ%`R}LVMiS zf|$SDS5N-aRoB4DZ6Ia(@3|2vNtnY`^i2*|gEQZo>X^Qn{NIEAhc_W&Io9>XST{@@ ztE0LR+@9~O)$L&CvSsVW;(r;r-IF=yExtAOEydWk7GvKwaqQ`;UuW!5!@K|gRMNTK zS;W6YehO<|XfURNucJaTA`XAnD+Uxir z@v~I+xa^@qfwW-&|9l zF2;SP828zU<4#|_SdY8bkX^e+dRhIykKNj{e<;RyVd5CGRIk->>&xQ*J!;SC+3Skx z{(ZR^>y?RPbyTl{+xu$f>d*DKUdaBJLCXuEW4=+0`DQWZTNB5ex_S@X{$8PV2i^Ke z?AI#wRC+#6;^*hqcW8$UZ+#SP)$m<-d#_GjeGG1Wm(V+uU(dAqe%PyTXY~=h=lhuH zs?Uwp`m8@_#DD%k%bAyJZ2Kz@V)v;&r7ykO-wT1w^>4M-zUKC+KCAig?h)PIJoTw& zufC&g&wnUa;{W(5x)`*s%vOA`TC%|H50_s0TOSODw(bo&tFMXQnp11M&gvUl^K9qT zS$zw)Tr{nD{DXEl-=)RNwq z=I25E`2UtZWXu>pLbTqC)xDp@zZJvrrqq#_ct)_k)^lF)=&mtcenybb-{1>2_zDfa zQiJzz@O66ej%p+Lp7mHe^s;qUJ2m*gUc93k4sU;7(pilLkLlWLX+pK~9^1>evl`!v zcT{`B+vj9QbpSYV4gr^O4)4{kvpTZDk8b-flW~u0*iUHiliL2v#{aa2{fu6`qdFU& zJg)^O&l|yIo;UaE*IC`#;J3H^C+2xq!~UCIyra4Yo^c-qXWXZI`E^#$Hu!T5{%V8& zsli|G#XG7u;Th}QUbfEa{RaP_!9Q>CFB<&IUc96F8lJJf?`7+(CgFrB&x1*O@s6r5 zJbkC`W#c>O2A{6MXK(O18+@(?pRd6eXz&Fae5nRsy1|#}#XG9y;F-(Hy=2ZSau|KDxp8Xz+a+e7^?Yzrhb{ z@FN=h$Ob>L!B1}RQyTo720yRC&u{R{8~lm}zp@wasIG=*Z`|0+)>-|k!EbKx`y2eh z27jm*@2DPuXRIfB**dGI8vN;AyrX))==*XnTW9rZgTL0`Z#Vcm4gPL}f7IZAZSYST z{EG(vs=>c*@b4RZ62AE9buDV}$s2si2A`_IXK3)58hqvkpQFL&YVaR4_CUD>V3u4Zd1~uioHmH2As=zFvc`-{6}z_~s41MT7sW!MAJh?HjzS z!3Q_W2N=2EV=+?^8Vj_kP|_4~g+aFCVT?4gO4nzue%j_Tqi2 zx8aV#rwtwkweMj+?B&y0eca%m^x}Q1FX64v6g;gQxJ&oQ_WfID^%Z*SLmtWd{g>QV z&-Y(^hC%B;YUI%FUG?eP;`)pnSU;I0Safcc&`0z1ZtM`;(I;#9bJ^+{-w!;kK5d&r z4jao6X1jvUxFCSE5tm3**^-Upp&M4mW)U-Z%A1`UrP<4>mE&MKyfeJ01&HFD&r z7&3-GuNXI`YvkD4kXTdJu?G4ebE0kX);y=eHkh055u-+q;!Cc4^qV-=)Yu34RYh${ z%xR+U(mklh){Z%CY&@GDw0o>QwVf{ZUiq|ZJ3Th@>78raHbZUe9yzp_d!N=kX2ix? zID2w$YcsVptLWHmZRVEdjC%5!1+87T>>_RLxy*`f;^*6erM;cgY}g%PT-UA>t*O!M zwT+$iFSa?bInRIBZ_e5_xNBTjI{KZ#Hxu{uTH}2%;+)ic4Pw9B`Wg5iwdQBu4mPCw zu6-{;X7F8y58QUtnB51A88v?N`fJ;#Z#5k{@5|%5$90ctd$;xJ(LFzEi>D#}oz-mU zJn=RxeDc=safrLos6k!B*B`kjje(4=ohLZ>n@8*1pp)7>rN;LGt8thHu4M0o`j(Hn6dNf>z#ZtkK|WHhBL!M!tU;Q1hPeU;0!Vsd+vO z9x}T9EShKR(e@oOW>jr({7ulu3{H(Yt4-m1jS25tZ3efpTqynDx|V&a&F#l;B=%`b z%&%{?1=`p_W4gxi%baqKTef3N{ElPm2H&O^pRC#eK4u6%02wo|^=PJ?$Bx)WbnV^C z;BT0ALNBfnt!&%t+P4~1+uX-7nD4M!Z|&jQpgyA3Tj%gZy|WsF-t*mBN3|ckyWTrT zfX5EwU0~dZuF>p3YkMx*pi$#Tw!atatgg_%-niOthL>x9M}z+wp0ip1dm8o!dhyAs zN8sepTCSdVk#;VR_44hko@nqVd-0CyIe2?+eXBRXiQ6-uj%sq=FxuxV=Nh;*FoNgz z?Altj>8$!`8#8K`A$6@*hnMd&o@f2_=Z8Z1U4`=LCQC%r+U6(ql8^C+>jh*xU6?mT-->S|xRqA)Ew{YL;A#mdNe0SQXdc4rv@%mKH z6k0oOpXvp))^)PpG|u@Av{pgtT76XX>(Tk{8NF4j`a)vgf3~!q`b@%&Nx2^;ZSXz~ zK3Ol`x0(i?`IY;62K4scovc~{JZ{vuuHo%6$E{VXAL&=W2doS(zr+8@#2B5`THtd3 zZ_)BCe}C3dZG+a{hn>|9;ITvYXkA&I_ipf>`gc?#;SAf#VNWp23+`JT0M4&Mt$X&> zA?PJP6kgWt@CHAP?Md0eIc zVCT_UT@Ow$<39|iwymE+W=@Z`wGmxocgLeo^%A`OTS-@>9p`1^v@Xc)`h4B$)2=)3 zpLKoNRNTn+KJR%3fL7o8v~rpeUe<6Xcv-_)8hpV9U#J)FTP*_rclB8k{a@?Tr&<=h z^?p>pA$YE?XpC0P_}!tHO=~UI2T!~otZT6eTAA-=4ZcN#@6g~oHTcdAzDt7-h5Os8 zti|w#ePn}=Zt%Ste4hs2ufdOQ@M9bN_y#|*!B2&IZ${l(>wQ(jeocd4*WkA`_#F*? zSA+kq!5?k##~b{~27eZAzGJ%DpTVc+16R4eGdB3l4L)y!&)?t+Hu#Ske3=GcuEBrY z;48!PTO!Z#)!=-znD~68KfJux*|gDr^9J9t!G|^YZVf)X!S`(Ny&L@K20ymJk8kj^ z8vL9FKd-?rXz+{SJ=Z*bL({Nd+2B_<_*)JBb}!zy`VgLb^Pb;=^E|glZ|(cfdiCq9 zzG(2T8vNUW4`5v<1vKqkk&Axt2G_KE2 z=L7o+pSMk0G4mZuIga14_-bm+vGps*Hg7e%+t=?GqM5In-vLD99i*lCxjmXS&~Nra zSy+ZT54YiF5%^Im~R%m{=p|(|_`JIN^&kD`&G1PV_ zH0!IjbD_;qYl92T?<@2h*3$URKnH+V)+byaU+zCO-P~_}Us3Kqx%d@%wT3>IZi}dZ^Eg7a>}@`%hgyHxb->`pnD!I)*%b^_`no?g@4K%nNrvsOjfe zuBTX^c{#saOMRmo+w;O)qv@~je6)@qUEhV_#){VV?OQF0R>sio9G0VP&%d*E9_X)T z{2$lY^Fv+#pU}4F+^6EFHZ5)}eYCGZn|N!2?WZ;yb6ty88(W~|+rQ@4X>G7}HOF3u zR?XPPu~r+<8c!d`+K|?GuA%2rKaNQ_jJGkZc6H-zLaSyxARmH}?LtiTw+(ebn3w2hbYRy&ylZ=Ej+uGvFXv`|77HcKeva!L-IZh{GHX zp-m2srQIBkqD>A*gYBc{n8(ms|Mr-jt>1R*tL_-b(dvIJhhrR1n=y=~-7!w3%^0VF z?W3mc^ul&VZBsMenY4*#KkfE8pEmI>02@zD+l7VgBCs|!<6TUfc=o%9!#)#e6YmPJ zebhWJuB1&LZN|B-p>@^z%?;gsuK(Q)-Myv$Ktms2>%VX4u9f~zHFWE%{!&A?PU^2W z^ov@0pXx8Q?mTC$wGU{gr$;tAiaoV-2l+jlCDx*uDqEeoF1H{jA!) zc40pk?fdV0;$Hyw-bw$2UW)_(gE{jok4rftT_3BbliR{|iix;x1O#_A+>4-3dOa_H{4a zqc5-6pNCIfbM5c8?W=nZdmryS=4MxU4=yeb?f=VB#`sUz@xAABT-WVeTFTJU+_Q4; z`_wX*@8Pbq*V#VQ!`Miz{(a#rZA-Vme6m_I{^T{bPvXslW~cajA88EjuFHyW{q4UB z*j%hf>;uuVPWlgqdrug9H{5&F@R4xOp&t4Pc!2Vvda!>BMd#^2b+&yWntD61f-fycVj^Di-C*1GKOYV2%CHMRB zlKZ`P$^9-`?mV2=&HPWtK|TC;@P38;arit1zmb!DW_UmPzYq4@_nyLi`8mXx=CA$B z+I}EBhw7TEsB)OOF&Grf3J2i>y-k=*Bg-f&&LMtiKl)x!AB!I$U2S13%g zRU_m1dkI@>X4`>t$$h?Ca-Zv#+~>I^_jzr}eO6oY`5S!U2KPB_{C#iq`E1F3HXClc zMuYooHufv9_b<5mfCl%uZ2WC~CR=i!#g^P>u_gCeY{`8F8*ctSe=Yfz1y}dEYwWf@ zcP+WkT}$q>){^^-wd8{ud{~3~95w#F*ZBN2+}7u(CHHx0$*(B5^YeLWY4>?)$$efL zuD{Pq!)<+DT5_L>mfUBe;pXcz(QsRziI&{wpCzBI!F>i=+WR&5as@YEpNqzSdAQF+ z!@Vc*d1$zHpNEFq`aCq;cpDd7f1iiOZtJtplKcF#DIbz9uj#zS^A(q_dh$Y{- z!F{e++I_BAa-S`RJAa=qhJOL~xnj7j&lSV<_qk%Y_3^o4_y%yFD~9XubH$STTru4C zq6YU_VrlmoV#$4uSaP2umfUBECHEO($$frUa-ScTTvBqM8HQUQpBaYR`phug`TEQ- z-1_*;Fx>k1%rM;6XND#BSz*b2Mp$y650>0#gyDXU^EqL-*q5odY| z{e4CluD{O+!_C)cgyH7vGs5tL;XWe_H=fT3!;R-N!f^e4Mi{QY&j`cypTFShJ|m1> zf1eSC>+dtdaP##UVYvDFj4<5RXM`p98DY5hr2Px-cs?VH{R+6x2*Zzq`;0JLf1eSC z>+dtda9f`dhU@P$!f^e4Mi{Q$XN2L}eMT5=>odY|{e4CluD{O+!@q$0j4<5#_>3^z z)@Ov_#`76r_!V%U5r&@u_ZeZh{%;jrf1eS?ZtF9`aKB&i*&j!Pd=d;0ZTb~Vv zYxmipd`&(x`2Egz-2bm>;XUXsU^lD;AV_P(bTh#76Yp##=>B=to0ILa|>S*tWVZy zDX>23S>qppv&Q9KUm9IoYOoC0{_2UpEZBM`{&Hab)Z?=}*ty2%$6)=`t)-uB)#ASr zSS{nO3|8}dJab+Jyeda>_IoD#`rXoMwQcjd=0DZO;mbOEj;{{)^RoNd@#Xr7*PwMg z@tU-bxiY={o@Xui;eYlGFw^JN{l`ex)g0n@r*HT(Ccm8buDV70RU z`f&Bsa09TK`Q?1u5U%cbn#LUfRugYT>p3OfxWJnfc+&#=`91!d*Vr0chpphw&sba1 z%EepLrq4EDzl%%lwuM_e{q@sspDk#$W$r%%H|PFyH1*7VJFuGHgQws2@bojTTtD#+ zw8?A78n4NkdY#`Ho_#bBT-Lk`uKm7xFYN+W^BghmPGGg1M?=7gH5i=zJQS=gadri( zCC(tQTH@>mPMmIV;_MFAmN>(~YKb!ptd_Zr0w>N$a5=ZpaP5h;2Usn!Mu63_ZMXdfyU#!N z0_$f#u{Gb1*7)|>zqZ*L^8nfdITG`r+HRYi4+b0KhC1g%z-sPQ$Cdwr!&r_hPTwQJ z_65`V1sgMa`#7+gbTou#FKevrEDkmE7iX+EOp7-7B!RppDac%&2 zaHKvrg7s5Rznj47K8K9oui)mT&Go#3RxR)6w}8v}IiBz5x1wvy`}u8PHD9^*-wwZ7 ztv+Yu`X%Q(!H(y+chJgL=A8BZ{Vwo=wD#AQy{4Z2zX97nbGaL==6!AQzZd@R^8YQm zw&Z^wSj|`RzaM^&TAjaKKXZA2*71_>gJAR3UzXMg)?caG1{KFje4hq)P7uD|C? z_PedN=V)`yd>-t%VCyyR1=<%meA)MpwcR#zc?oQcTpL~ntK}Lm{{x4yj4e*zSHRXT zeP0Eu&BWpP`X^dnUdOb(#-U~$apJrQF4y2KxNBfNj;8%H#~U2_IcK>(&MEP1UuPW8 zBagS?iT@XHnddujwdDC8IC*M&mqX1w#fkGixXkkdxV3wLJUu@?1iR1GuS0VWd<3?i zHf!=%T3^;g+s7PgjwQ}mpMcA;K82TK{SEF|>W9?he+IUnw$$!(u(i{c*nbDBhkpUK z-Z^K#1gq(1y|t-X@8oJLPn~R=_1EtHb8fa7?<=r#$#`Fb)$%^BP0jw<(=Wi&{~NHi z$-es*td{ls4xAcm`v-@bd59C|pJ3yJe-GZQ;FPsjT_%C+qn`Rs3bww+vQ~2a%spe; zHfxjqoppboZzqpFV0GhdPG5O!eZkYArFN6S)ttNj+SDBXY_M~-P5&vdZBpIj@r-GYv)|FKj83XEq_?s zZ8NXA!N$lu<^ika+9{ut!&v4ZPTzUK<`O&GmK;KL-19 z4%(LI_z8!3h!bZ8uyInu6~Ss{4OfC&Lv87|GFaWY`#e}~u8!*%juqckz+QLaD=&TR z6W>+A#tUB!?3yRPpK9aK&v@Ee8uRo$Y7Mya@%rPv#+qR7HPnsS`tC+Mhy8tay_OvR z$>_f}Sk3#d-fy#(-hOtjO`id@>KT6vu-C%ObxXKDO! z`tfKda+trZbsUE7bs~9m!^``G-QeY#?GD#RJ$;6Qoquv40oPC6>(3}!U*@50B!`-L zh!bZIaQ3A8Z4BJ~rXHWMV0CjJO)HOWZ}1@ch3^B`C;M$*us-VfPHjK%UL0;R?diKe zSX*i&x8}*)Hom_ATSw=55UsCo`NX4sAjjJ918CQ)+nE_l{tBPhZTbIJyI|c~7yCK( z3$%IPHnukVWb8x0I~8+26z-h!neH&KpXbz5hr_}4)8-fl)2b!rkzlpd=qRw?@%g=h zw)h%(gggnX zp3mB+fXgwB?{@{KqHD`%?bE<&zHC8(6!}ybw1d>>W+67tykZ<>0y4f3Z)1tC_d{SAgxK?zorGs%6}( zz!}&3jKsYfU0dQ_16GUuTJSu@xYxneGVb+Y`=~qam9%R1H?JGP<}x)`R{2d}hHw4t zl{SgLf4A=qaP8@PGuU|PdkZ*ywUvE;1=pVW-UiNm%YAV>y0*;s4zOD0dndT8^IdSY z%=gz|`=}?c-++^s_A;+q;o6heJ>cZ!@Afn9z3AGK*KfgUvEK(S*WrG+TJm}TY#;UH z^&mKTX)p7-8?HU^9|mV%yH1b5&0V|vcVO2i;D0ka2fA8xcegUo(HQX-XFlvaM;gya(y!A z7r|u>{s=E?@Df}f^^EZ{SUoXb0bA?vSHb3#I=%+hM?EqA1Xj;=`gO2z)y?w-TD9z{ zH^J^HKgYPnsmEJr+A{W^!JbbU=WVcB{Qm+zizDN{1J*}9G2R7V!jTy7fz|T<@jlpo z>WT9KSl!RC&fC56A>915Ij1*h)soA{;H;(oe}(IlwfhA85r=;Em+O<7{0(ePtVenN ze}<+lexDb9<@NaQ=-QI+7hwCUr|*~Gj8UG;U!iM@-`9m-dF}lMU0cTZ7HnVj)b1Z( z$JQ3@JFvR-%oyLxIhuRGVr#QT|D=_R{k^Jjthe{P9dN$~(dKVg<@sK@6YTfO_P3vQ z{XT_j%bN8CyJp(_y{lZCpZ_KY>#xnENCV*%Q;l{o4rr97nE?*x%RY{%%IFv2FccW@g%1ID9$wthL=XW6uUQM((|52dm|K z8Tm{c#xl0}KYc!%16w(dIpM~BpguR|0=rMt^L@z=!1mMT{qNkgzP#7h_CpRe#}+5= zdBEjZ{ov(T^THiVJ^OGzu>G{9w)2Cnt+vEo0IVLqAlUWDwO}E*ntrZ?XOFm^&HPf6C9yev zN3lnjf~%X4zu}f={H4LpN1MOrmOIumwB~B-eq4@rc@AINe_Y#b6Ke&qF~U~_J6`tZ zN??7|-6Qg4IgIUCV)O8K_)|^(O{V7LsTD`X0CWke^#z+opg4J@4 z$XDesmN~RAv7uFe{>NMrYi+Qx?yJvU>xUA87aDB4R z*9WV)|FfnWz+F>qiL)VCJ?G*^U}LL$E^b2W%Q|b@m_yAuiW6s3u;Zuha@RM$wsPxi zU)%KE3~W6z_szj-*}K}*?4Px>&gs7;*!pIkTY)#Hp6c=08mw-9ThPj5+ZJ4Y=KmR7 zpXBp%us-U(C;6L|?ZA66hVQ}J6Ki|0Hpkh9Rvz1q;Bw7&f}3}~a~%j)bDvuKUBJGq zy|ylnT{*0uIB^Gojgz$>3|1@GehA#P*Oq=m!RqF^Gp*bjJFa6mR%|(UVzW)0VPN}b z{BE#wRk!~Rv}*C+9c;c?ui;?trPbpz0<3QT-Du^pjRKebN5j2$R*%mfVD<1ZVDktc z3pQ?kKRFJpk9yX9JlHte%wr_2T4L-4F7wzM?mf49eD(pWhwlq6^Vkn=9{COC{$PF7 zlgBT>#?fXTd(x_j51{pLj*aiV&Vg|6b=*G(gMGPw)DPlNb8W@R=`gVKFy5iG^29p= zoO`Fk!E*0yyze;@{4uS4v?bnA`fy}@j|O``s~(?Y!0O@0g3CU~!7reXdVG!ttB0Qe zwidZ}JQ1vqdVEd-yKbrF$zc7|)9)0py4Ujfohs+h=66N?ndfO>b?bCGt@9V3L95NW z_-_cD3ErGmA8q;^LaUy<&IWJ8@o@bfa1L0_bM|U9*YaGrdfvm%0;?JCJlb;I&xdPw z-7lo|<+`h1z@g^4i<85};PQLFOW<2@q%OY%`@Nrf>T)UAe%j3AB3iY?yd10+{|R6< zKd0o{c?CS*1!_;Gv-yaYKeIR zSS@+o2+lQKTl{VUYct06wDQ<~1y;+N-2zs-nL~fqNzLnO*1*;={f_TeuzSPz5s$>Z z4NY6(-VRoCFI>U8xu0Ik0}}?i-(L zJP$sC*7(}f_YYum_>|p6Q@z0d=xdJd9qraM^|6&NPiwtw6YE9rk2sRYAHiz*?SwWp z`!7Rl|K(_#pDC=%OX!ZD*ss9-ypZ|43U@y4Ic@QK4Xn)=FVo5$?@zRjXX`!u8?iaoIUoK4Rx^M1+q+<2#?kf;hnjK3iSr)V zd8dBwgVl0ht4%HbAAr>|?uTGC-FR-{ESh(U;)ls_hF7HFFi4tLr6Km#2od8Rr|Yx#xWO7Odtu z8=rrG&1GtKyw|huz${g@K8Ln^pZ8C&wyekZVB@Jf$FFJC5_b|V&}zAlnH28l3ibGO zfYt5q1zjFnC)m7lU)%?-=D8T3zO_%TF_Xd7^iAE>;y(qr{H!u1+;KCfslfWE=Y4Hz zaNgInC)PAzZI0vL?8;-C7Q8WgB71QubdWsWgX0SDM4FB#| zuFdmtR&eeu{TpC;?k#5ndv9qUZHYHKxLlJt;NDxR$7fEkdiY%6vd<6T-dn22=Z9eR z@VUVYGym+ndEok}$EP27etc5F_43|wez-Ym%f00SV0G)XAg%KkFGQ;? z_m&HTy|>gyn?C-1wR-Ye6kOidEe3awB&Wr}-q)!o_a(sg(`HQn##${gmjbKB|3_f8 z^1g0qc<$@8r|&XgZN^=aRvz1O;4<#=@RnA8|LbS%AEPC%_Vir=tj)Oo&9^+ZmB8iw z;>vKx&i&#~z-s0F;wo_aX*1@Ev}%dD8h8Zj9sX0e<0j_nV6`&l8gToyY{c|$&s!Sj z0H;N@Cfsw+>-&1NzI^VlE?*nI7Ol_w`_uZoe{F5(>(H`|t95A?q#Z!(`KNz!+yLyD zxfW~)PA#>?ZzHfaua*8gHyeZ1ti$@Wa((^Ww<&l9>Y!~Cuw2^&Y@37o*S5{T@{F-1 zctCCQZ|LQI5AU^LEAZ#E##xmmEKGm)A1&pL3{LOR?`wu7zAZdF=>J9d-cAlh;6S>aY`7uFcPfUEmox%xh;_xi;%Q z2yDHz?LsSGnF_J&)nM>#>nRq=QS2hJ=djiU^TC|={FvpevT>EuUuEpt;}O@aOSZWSe|+83(h?D0n4?S z%l_c(@BP5??C)QI-QUKupZ1J*09f7J99N$44g~A3&2i-!?;x;ov^lOk`}<(9W4gZ& zrS;|hRzHM8&HXGc>!6mTSjRU^Vk{FU$3Hj~oL|ZI1@a zwN1cw95}T-7A)81IdcLy=gje7xi-(4lfaJScqh`zbIzO$_MCAnW7=1rjCo3}r@nG+ zuIFjs@|-yxo^wW9{LTPt%YHi(d@6_g)VOke;(r#{x@OOw4Oa8MFtt7ho?2^9-*ds* zQtR`;YL4x{m2f`TvGW;2u21@304~SA5Uy6fk6#4O*xJ+gVz9Q1eF<1CWB(HD*uMWc zZ@E6{e<`>e`!cwi&zCdr%i$SYd-_fQYs=VIfYmbgm0-tSj988>*T??uyQ{(3cUOVs z+9qJT7My)|4Op&i9JcGh#x>@3wDMf5ZU8@mkMZoMJ$v9LurbSP)vwUhv+g&8)o$cS zzgxiR=a_Q+%5&8B|IF)laOQOzSe|*^3C_Ik0L!(xw|@=JwdyXgJlCq->SABURPDsfo{^~8S= z?B2>gdI+3-q%D3AgSBNJJpxuUFRvkTecdy^2dB2b1Ix8d!1fq8wS5#U*XFhA3Gj>@ z#(bPsuFY%JQ((u@_9U%5*Q%$%UaK6-nD*5tV?I;s_H}HzHrMkx@QNIn|MTEntF*=M z4`6NCZ!dtK<;b;4u21}51efRRAK_}{wdy5!YOOteUj}One+68|dKIo#UVC1HCzkf~ z{S#PQ`0L;@)*EoO@>=pHJh8N=?^|GP;eQ5~vEGKO`57to{0ls>w5RVoU~QXKlePXm z^LN3{+k72IyL-)lgYrGFc5DAW?FSszKG$lwK8gJyxXkM#xSIdwZSwjUp1icD?_a^% z!ao6*u|9>X`S0Z>*5BZXr9FK=18duiyg#S?JBRVrotyuD<`=c@+`a^V#o^rYULe;e zF~0_v`FsOc^WW=DKHtKVkM{Ka2UuJ9ci=MCKjCWad%*Vd+VA0sr9FN9@S!bpn*?rr z^~`NjuzKd!0d{VAuaWDMn4RDVEB}2lxwZ+|767MK^MmCXV*!OMY*-v}MTLi3bE{-eDc#DGd z*XFqLjJFurINBUnp7-s=!H((s_L8)|eBV}If36Ut37?!1Zy*{|0bL~wzb6^ z4{&an>pEaHe}C?N^n3Aj;p+F}BOd@((T?K~vA#YztQNZ%ESbXYllMOu2sL{$7{$^;$twZHLW$>Nw?C z+oNmC_v<@=)$;xNPGDc2Puh0mQ1g5en~%?P2ZD`XKF{45O+B%@z-o?fo{3|dXSKV4 z9mm}DcMQ+#!L+{2L!Ut$yK7xGl%`r)U#H<0IRj{ciZo+2f)+MxytoRj03?tF#p6j2(FfU*_N&FKN!3t z{^@@RT%Y`n=Am#;kwrQ5)$X|V(N^vo_fO&+1}^vA;c&I&p-ru<_W<^I`X5>M&ojEC z;A&-mHT}IV9Rtq#91WIhn}F>&aMtHousq|O0M0%=9xTs3JrV3aHJ<&nXS|cZ>ej+> zRT{SDx`s0UJk~T=Jr{c11>dEVDaO!XtSf0Gj z1*Z<@faThJ-##Cl_wDn*a&x?Z_DqhvPhSZ3ecIT@&{v;~aZ#;j?s9Fu_gw-m-=}{G z&-37N@VCEy+H>As3pPgio^c(Tdd|7)!D{7q-8aC~&oSltWnMReJ@2yLZ-T33ziU$~ z*L^+C7x@gd?u%QnC*G~#$}S{zJqn|DO6`xLSUfqfIUTkJSF*zk~a? z<@wF`@8SBW$LCRS`M2ed!TsCv^m!bvkNN~wnWlPz|1p*}YyFfWFKeysNe(q@Eq1

2rb@;z&=Uvhf}U;mCV{8{+?@SGRV!SzwkXSV0T`MX{}_onY3(6!xA&*25I z{neerA89iOZ7*`DnX@=?UIP0!jNvcC7cO#o1+I^Ja(We)BgixwZ+|{sPYR>}{|-_Z{zm zuffMY+A@cC!M=Crx9#u2)x5_ruMfb!+#A~7=TI{*apHUgHco!K@G)G?-^)e&D_p(Q zK7p%y{~5ne;p*n#dHx|-En|HSR@3$wt$bznh5uckzk_F{RoDJES~dGxzc0aCkf*jU z!17$nzpicNwfq}2b^X7hmFw@n7yS>g@#Am&@8IhCe@iRR`^Z1R%i!a<#?tOSNj+`l zS$Es-2)+`B|GgUH`!{gDrs7zUF}OT&Z`OmSsj>fF|Fj&glmES^{G0Xuw}$MKT24ZZ z?VlP<3fIRuIS08usb5EJ|0i+9$<5kZcVGUuG?Uv*U~`y-Hvb)+zVtH(x&LnWj0IPp zx54Kxxbs@D=H~40PNwGYd!DR=_ZHK@9s7Mol22dzxzBvh(C+-qMO*6O?>E%KX8{{0 z*Q{CL`lzqX_{N?M?D*c_=kxgNaCP^zpW){OtNDGPKK7dnte$i42Vmo?J2&V0L$LM9 zdzM_k*yk?n=FFC<=0Q`>ePlndTIMt#*f~0XYb}p$0kHGcX5HmtYi=C(vioo$uzS(| z`e`>O&x?h@=InadSFV3@UId&uCFe!a)V<~`PV38Sj{0I8YVISkaa{K$IIO$1`D?w- zOV)bU+14@KpG$!q*S#w?o^fMey0DA0x2>uF4KiP|ak#gg=j=6h&-&jiGoJtbT;nfS z^mkwMP>mSDxJ0 z1v`c|Mn zDOj8NIWKu^n}O9FL!3SBp7Gz}_O&oa>bgjct(~9o)y0c(nBNk#o_}gv5>J1Yf8uys zfz>nq)?nl0`ne5UA9X*ona8$Zb#oLaN6#DIFMTb^ksOz*u{r#x(Eabhnd7pw*#|oi zPd~ZsP7OY==E-+ExH0ov!R_H{?o)ltcSo=}Sa0LY#m3e@_1PJ0PGx<%(A4#JY#;oT3N?0qD{@$am1&d5AmZp#<}nyeJ@Xy{wr}pi zc7>}A<;dC%1KUrVYiB;)V0G6{oVD}cCi^LeFaLYgS-VwhY~HJJnCt4aS)bv=DA#Lr z!M)~=Ex3N;3%(b4@0w@sBjAox&V3}Bde(RpSj{zdUCco(wb=u#mO6|9t4+s|JuwdK z-0W|Ta0uG-8B$r4c5lC z9*3`WII;%o)_6_Y0UWNu`m|Yt1Bg+s!65~A4Gu53en&R=(FNE4xPt3{LW7@FaQ#m! zxc+Ap{4DU4T;qe#)U(D1gVnOe&QUEjKNPH%H9iciR<7|8aBJrpSOdA( z`RkuGJ__s_>!VE{Yp9;_jschB9Sc{Rw)XKlb38cXSwp$8te^fF??kZU>7z{_*FZhv zoeVC=I~A^WO6}u#r-L(|Yaln4_18b+oe6e4eYEN08mec!bHL6ye=l(!SnXVn_?!>c zCu?*eSnUFi_*?{bE{St7Sgnk6N$rz3zl5t9C%^Z&6l@NTVGMoTKbO&7&f%V_IGl@I z|KupVtz@^CJAlf_whlTyXtvZSdO*uK!&H*Z(&K{{ntb&2zq82RAS4l^R_S zR?iyT09MQSX3lEyzX`0C^X*q)wdpwYaX;PywiedZImyM&NB``{+raKeeYELgP1G~q z9pG}jJK<{6);^B+YjDQ1CURq05B)RV-C)PlN1Hy@R6XO}3pNM$%!c^>7OZ9s@1wn+ z!@6eA$@NM92f)V6=g9}bYW9DK_F)eDXV1y?vA_H3k=pK_6KBtPPqhPwub**b&;7i{ zThVS;<85iT=Wze+NSpokDDkrY+&@njT>aSwf3D#Azglqp|J2~G*F5{}F}OKdtJLRl zuzL2}6JWLMH*-{r|C3;~?6;@DYUO@=25uc(Pv;~TJ0Ja%&+}mGQs(mqGxQliOhkIvN+U%V-iCONQ zcMGomeuIBdaQDvV1=s(J2LH0=S*y3;=2X$9=6?pOXYafXR?FToN45C>1+136^A1?8 z+&k~Vt%G%SPI9sH(Leco2(~U|J|Cf}>+gK!S(A^!{cty?v9u?zzk2{BWt`n?I;fS&K|VcJ6{vC+&kYFTzwKYXzaF=);zWU27kvWYyT~pdiKIUz-rkG zj;WTKe+O2}Uic?it=tQwT{$o7WR7yN`RSj0I^fQ;%%>AgU4QeJr^bE2{cty?v9vd9 zyszf*0uyrrHd|Qn|s$W#O_=BCC_ES=BdAa z+MDP45!ihl!jW_R&>H`O_OKcsOnW$o=lYSfIoFpXrst5{b8F>-xAX>Iwcz@%Q*iwU zH28W2*MHN3>%Vz}Z&7gle_n9?w{P$r3ahDDY*Xo zHTeDo*Z;7B>wiRpA6anyPb|3pCpY*h1=s(ag6n@?gP&h;{Vy-L{#P{kl{L?`YI(SO zsH)G&zQ7-&sps0c0$45APWO~r&Z!l_YPoi<1Xe4roj-xQZ#@s(gL3gIwE8EX)xhrK zGM}HKsq63ll;<2;9o!FhV;W0)@>&C2=CvkVt$pq6+qx#L1$R8Joyphpy+69PobPLc z)pEYC3(kBzN95+``QbQX&kg$}*Y&~Xs=t2Plj{cHGS>~^YVGT1JJ*fi$<;a7*IYM7 z*Opv20jnj~&A`dk^HXkqo|}#%_Pn%Ta@`VauKMezJ-KcLE_2-)uC@q=&oR8#Zv#%Q z&LO#Oi>@uX{tT>^T(<)!SFa6n^YgmkIN}{>?U!741e>e=`e{$DJAunw2g22cl2dZs z8J=97Lvrl`YfG-XfYp-gU~qEvx+FJ0uSt$09!hJ!ORgipYRPpJIJtUFl$)Q|L&p(&Ewo>99RoI3{q@tHT*rdTT*t%J z#&IOqJ;BMf+^>6~YfG+sgVmDjzTo7#4_Iz~`_Vd%cz;^^CD&ho%~gN>v?tdCz-6un z!PO4rNUjHi&2=%_vaW}~wI$a>!D`9%aBy-x3@kUlBWN8*?0bOylIu}mbJbrz?aB3M zaGC3|aJ6GNlIwBcM1*<2|%fM>Ma{@SdUJjO<+ZD8qA-aBAo}%dLrP?A*kzuYFUWhr#_ALx26WH}~wh*nOSJkv)4> zjZdRJyT)hGp2OjuJ&!hf_7P&1d-jQft3TD?PuDza{yVs1lxzNbH1+JkN5N{@gN~_| zwSEk&mOc15SgqWHPr{v-buve}*!=WQKF@%iXPM8lXzKc#zdSX54%`oSV;W0)v&NSh zkJrTzJT_ZHNJ@UQVwf;Ic;kEJTc1}zg%$jR~!7bnrBXbfICJxrx(!FQ{xxG zYN@efs-?z%1goXSFM-v{8ovU+gkh|cIm*T6r+@PK6WDo{`Mi#%uD|)qQ{y+l{cty? zv9vd9e2wvVUCEIeUsdBPXs_n5uGiA$`ua9;Qdiq|8vNawC+|1m#wqiD3r#(B{4-e1 zI_BK{3*20sn`6qwj;Fuav8-k0{2sg?ee~B)dvne=VfS?%N9KHejh*uiHSVIlk;D1_ ziZ=8Ah?wR4|JvZ66x{chFADDSgs%$j7+=>s^?Dy}4rRSQKvPe>J_M_k_4*j@Jgkd3 z%EjiVzu4To&Zb_U!u!!jfBm#4kH3MNd3=VZp8fneSk2EDJ{R1HSbvABXC1!;CogLu zHy7)muYNnz?;Ei7)@DuQ$@5$AP;Ba3jUGJdH zx_(d0a$P6s=*=fBxa&H3!F|R!Wx*X|s+y;s{~*3Ol|nLBUUR~ohjlSWx!C;l7n@smQLi7u`_V^#{j?{Kxxx2hEAyBKO+9!7cGyU}kUu=UnvP2^eEg~6k-+24NJo9p@rc3%&0 zWL+PuvFrL!jmOhI%;CEJjyCJMI5Eq0U9!QKD!A*qe8GKIwnD)jW5t@Mo{PZEp{(bk zXzE$l#lULix-J2C9@fPil4_0{hlN1`e=<^ z*T-snFzw?UuIrPuS=ZHwS+48q4ZcRfUDtIB?sJj#3ho%|*F5$73EUjYbzKEbJ?pwE zSgl;wpTeDobumY|*!=Vto7;iKx~>WDM<4z5)1Exm0yp#MkEWh=T^p>Hb={v>>%i5s zt^>fy%Ua0I#X9J#-$C@-0BpUrSrd8IbwltG*z9jV?ag(44!f_XII^x!*VuJ^rpCw7 zKFi^{K2Mu<-ISQ+x^CX!TNK=N{aL|}hHqDJ$JoB+spm#;b12t!V>I=w>n32ea$Prr zI}ht(j&ia2=`S|7V~Ta%65fwK`s=4Xd29u4=CL)Jde(Ituv*sjC}M33SI@fs9GtwY zh1^`MgTDG5OTQh!)?1r3k!M|Z1fPt}{`S+}T-O(|`zr4#{s4E4UDp>lT+2VwW-Ysj zQ?BKp1|MAW)MO{PamuwBh^C&k*%_=>uFWoRb8&8tDHl7Q{$j^Ey;z$e@P72sUq9{5 zIlqS8*GnAP%P-g1IloflGiYDsaQ=Uy&HTHGS-t8GUDr2jd=c$i9Iorzv{~2viCM1e0S$g&!ClwG3-0$YM;81d_)#@aJ@-r10^RO=FC>NWb{$g|6k0bRu2;7f8`s=4Xc^nLG=5Yv`de-$&uv*sj zd}192SI@d00Zv}lLT)bBL0|nYq~Fotw`jFl6M6DH27DPd``b@@b6wxV?&~icuB&VG zPK{lMcR5_k_i4+%Q^G&y@bv*l`hHmBDQG|9ua8kkTb8^k|H=@VF zr<|nqQ?SV}b9GZImM)Y{Fn$JV>UG52Rb?f4oet#o&JZ;u*a$=nVHcxGiEw_Ho z?^N(~wDz~3c7GSFp5Fx<<8uyQpK#>d__W5J8-J^@^Zty(n*5#Cn#@36=M%52$%O^K z2)_vhcTKJ?xbd$oxc=AIJbUsCxH**P#+hj9)}b$^v%uz-zsWxvu8;bc)aG1TU;Ztg zwsSbtoQv4HEXQ2r>ekbFT)>gPF~0~rGcnY)J8w1fbnPz&tK~OW*7g#(x;b1*>&qO} zf61Zdm|}Cdtmf)||8HHd02|l&mo>N&O`G2-c`jWARx>wiF3-5yr?lW4)dzxUrrpxc<)+T>qB~uK%kA z-xmIM!S(;J;QD`DaQ#22xofpG{BF2;-%gz-r@a@hb`OVp`94}-?qzMi<@hs)^AunP7T z^c?sme2iy5?b*|R02{MBCtg5P&z^n}toA%d`u!1{evT>E&wpdWx>`ePnR&g^vbFPb zU&%AC*T9*Vdra=%!+1XY3H)nX`)JGfuY)tbaXio7KvPfNZ-Ujlhe*G-;OS@Ha{bD@ z&0p^Ra~|%w?1jI8|J`1A2VI-FxVPT}ySEQ#-R1AY-{nwuFKmPM0oZ=p%*FFbEj9ZH zoW0<=B~Q)%3Qo;D&*a*S`zhF1#`K(%XD|E>?DzA=v!C|V<}+|=V=Qa>IhuNE`ggEe Vxfi~Gr=MfW^($*=Epu=5{{g6=_aguR literal 59320 zcmbWA2Y@A2wY3Ynd&oKGCJ7=SAd+*KA?J)Z^vv`?Lr-7=IY-GkBRPXek|a5+h=78i zs9*w7k)R?X@_*mGwYpE&)93rWQl7Qe-e;eE_BmCz?yc^gj%k;iuBxV~W~`>KUY)*b zz2>NyIvIJyRwGAjyX}arHXS*7;>3}YrySDXJF>rTeDBCHeUo}e zPw6|bSHI~CzrOLK`=^d69lG!+bCA!-7#sHYPZ&)*j4b88CX5|Bxpzv1G9CF>)%f0n zM)v5p?9`Rl=D$%ZuZ^vjhwK*d{}FqTekx3&^vi* zKO?lqZOtFtG8LM&7~9j=-#fsXBOdY-NNJd_E@{~za zM^9N{Lz7=&>dLL`+x_OpFXNf(0`S3k&s25851Q0BrT5>)p%D1kcM;k$)}rviSRK`3 z@B@4L@tv_+8a%1DXUxA1i`|$@(v~rof)B==srmwZa)00G|D0b(wXCr?jQjZ^aXPEz zj5Dxi!>SeF8}*IfeEj4oJ>y6B?$A57y+-Q326LFTS_NMCj_eyhruX2~yf6@Cp!i+v zZBLR(1U8?Q!L5F@uJNivW3p?8wKU|ax4(B>@AxSr`wrH3&T5UJz9X$`h96w=dJp`I ze2i&6tI@XW(^0JsPHt;~x9gd*uSM|ukoy{}L1(otd?&VOeB{1LpY`CzQ8$j<*Qikw zn8~2#exE|D$+TUZJsZ*5?|+|IonfoO?fGm#zszStaC<&o)uv!`-LZEc({1IV=Bvzi z1YAFL^OgHDH(S?4&3s4ojNX6DqzMz%&3Q{&`~L4a>p==d&YROOId1`O=RB<18ocqi zF*{C~G-3btZfxhI?rSjbu4+5@fG^lQ)qI&dIWg-&-Te3M=^s1tpuRD#PFFQwwF}xn zjFGJ?z`xy61AduGz10SE=C{IFt~$hYpY2GSIqn2*&#|N08Qfm;u4*?h{fGDU_Y<`} zS2bUQbM375fY*MldCGk?=c)cb*J++=Z}htVf1YKXO>4OZ=cwuo`<$Wk+l!c)-&etz z-$-!gH>y4E#uNG{Omc$mFEw9-^Xsa(mfX@>bJ>TOnTscG=F$gl&t+KE-yV0XzVRa_^^9pPS8G1%zRLOd#!rs5Z+l%^ zHv{Oo{#gq(*Q@pZ&eb8B_X92MkU@>xh_=QXco*+rzM&3|@t-R(n0TWWwZX+5TmR=K zc8w2UjMRnl57fn)q%H@x^Vp$h%+R{1`zq^lD4d*Y-}ccmOW(s7|QC&vN%;D?c z%;9oydk&q|Rqb(G3u69qUxWG2R$UJ#w~>_TKj%iIBw-HM(llY<80Ue3aXPARgWENnzk0kLvwpky zZ$|w$?=L{W8iiU3s%p9UB~+J``_g7AFkZ(T%KZ_?BA!0T%H-oWuEH! zI+xa^@xRIBUtCkaSB(2yG4A&V#+|)-sUCNu-hKMUds+RzkKNj{FBD_EI55UM)vI;f z`m*?ck2-jI4!NSbe_t-fdSzg&j_SwY_L|OJ{k$I63)z1&Xn6s2%%2uxzFv&^vw<;Z zu6_k>f8KB1LAO47`*lY>mBEjb`1!f@P1@f6)<@A+4Sxx5?~m!Lx52H?_(MbawMnb* zZ-?~ltlokTeor!6^{%n}%3|c`asU1UExSI~*!EWu#O|toPhWbqzpnwC>mO>Zeg1S+ zf2?_b-?%<+p1P{}tB+{g^B>EV_<#GzlR@js?8FDF6$;$`aOtJL^}%3l>)xQV`jGgo zIkm>?to}@Ep6#4EtG~c47fov(f2Hl`JFfbp8Cpm6FJnpYjI<>>RNxF6&=@>-aBcMA0b-r#rB`Jt+!%0UZ;+{#505SwVn%sC-zM08CUZ~ z8ho(^U#Y=YY4CL$eElK3qZ$tP-q}3%0=K{a=&X9clX?zX2~#VtF+=)wR%3_oj;asd zKF2z$3E;$;3@+mwIHX@^b#Q|p+Vag4pTW7U=gRjuwt2Ou< z4ZdcBZ_wZyHTcF2zD0v?)!rrQecQyEo4L);&&(h#?HTXOYK5v6B+~A8g z_+kydbb~L`;LA4niVeO}gRk7+t2g+X4Zc=`Z_wZyHTcF2zIlUh+2C6>_zn%eQ-kl^ z;CnRqR~me;1|Qwvy$wFL!N)cDga)73;0HGN!3};$gCE`C$29n{4SojP?~S;JZr%Hx z)v#X#-;d8vx#zm9VZXe=uN=a=s{7&IpZjSbF&-J>!}Y1bpJ?#s8vOYoysLTz?ihS} z;NedDUh-!{d^)RNH29lC_^|42cKos_f9tG%hu-?oM)H22CHFP>eO71n z4%)^O#*gje%YHu_$bF3;SwCqcSafc0&?oXVZgQ{B28Qu$Tz@VzgX8-Frq!oyGu*zt zef#XoBg^q4+s4G>6HNPAP5q%LdKWs=82J1aeHi-0DWm&i$oSK#x3h|AAm{0^^^6}s zA%?`70e#A(p7E1wLt@QX#~SH_$pPE+t$EIbZ45Wx<0g!sz?W9}C^s`&wtf#!P+*aXG%{xw!Q-x*w-6P9gB~1`hC5(N3-9NP&%qV)|#Jb z{|t8DweLlSwSGorpZ<}%PMEa+h)ENsPTYJG`>+eyxV(=}>6_9wq3zw)r$P7ps4bp` z`14aK=sdwTEPVRb&%F?Ls|llf`ZphcAdP{HuDz?88GrNeY#B)>wb@IJ&-cci+qeTh zoz?u|E&H%)Vfns&liPt>!?$8BgYGpa8`xN1Kr8PJR&MZB8hq6{M!rv3qvnI(r*u{8 zs(C()>7Ceqwmhs_FWRBKlP1&#$6p_P(wNk!v)TZD(4_ET)rN2@%Z1YauWQ*=ZDc=w z=HZaG#QdhKMzl135M$EF)+3T~{LQe9>p6Ic!9JU#7w2|sl=gZItG20a?!Oq^2Rqh! zYwy(t_1$W{b?y!5oz-6GgWqv=RAb>jOLct@08ieRtM`;~JrmhE*5pLA(G#YQZ+}14 zS)HqYz2UT90Wa6=+6KP?p0iN@n;Z6T4&l>PcfrY_Rp!B~e><0Zhxm3@_c!HJFA(rO`0&Ox31N4@bdk^`mLxx zKd_UZZ*ab=)LMHTJF73j%X)pe!PjW;wTAGnY8^Q97$}W-?uf>|AcL-IPw>#6&*`f6 zM%!_8@A%&KWvF!>Qu8%goz^&`YR&!A4=(oNsGiBa_;gkWpnLsvtIMa>+*tC%TYT_4 zxnmmqI5_tiePep{A6A_#-?x2r7}Pte)6vJY)}^C5Pu#kmbX1pu59XUT^SvB=NR6*j z=Nl&Vo77u)SakY+kw$L*>fM{8Z*>P_PupF?XEq&4SZ)vHCn zLA|Sb9lcen`XXoF-?g;C`uw56|JdN~HTVZZIPVYe%&*+n|3Gi=-RY|Nc)OS~VMc)>g0U3gB}8uiNr1e-5mp+7PY14?C+7;K{uQw62KGduRAy z{X44N;qCXLj;aUD^0xLvA2>fXY2A_0JF9W%B_9tj&!dSAesqH$)8NO!c~@%PSa{x? z)aH4d-OA-uwEuYy46DvUZ=E-Jvv1|tQJt^qiAMr+eDkxSq-T zp)1mk^R#hV7vy$*e%90b$z+7DDoV8dLJa8XZ z*J4eyGT*fueBA~g(cqgm_?8X6b%SpY_h+247P~g=yEpir4L+*D$29oZ2A|U42R8V@ z4Sr~Y9|`xKin_Jd`}~Ie!Un&%!LM%cYa9If2EV((?`iP+8vKC)&`%w!53`sg&Ta)2Jde0rQrFQ9?z#`;d~?*c&@P`yu8<0tI>a* z24AnicWm&T8+_LW?`iPS4L+s84{Y#*8~nHiKcT@-YVeaA{51IBHLtTer(r*@!7pg= z=NtTmf{&oy)9~S+OBSfAHe-kTtkBnz?PfqzRxADden5?o)d;RT_T{4%n4iZ7F9G%y zzIdCqV&*%Ga-4ie(bSq_=evoFZQg49*UHDw>(R_tKR;7P<88d9`8hb6HPCPVLgUnE zX$u#cbJefA&}OK$r3%gO8}wVY(EOf3ZTUhoS2b?A26FLx1+`TR&F>S`)+jW;Lr`0% z(442*28HH#25OrWn%@ydC1A}yE%My!`+V_>T}~oh?egDQCDb{Ck&M((e-{{8nyfD{j`s=#{t>Z`6cNw^`qUmd{E6}#**;%a!R#!957i%0{ zf6s~b+`6hSgXPB5NBdf|iN}wmwBo7F$DBM*w6WAJ-*szlJ=OzjS99$3Y1NEv9BbnF zV}ANL)~2+^bFGHcc5_U_VZ0Hv+SQG>8LgV}jAQ?;Y3*<9ZD@_HHZQhqTY5W|=bnCT zt+U!5tgpJY-GMgiU>xnn_gqZ;-NE)zGrv7*oB28Zp0wK4jp;d?nD*6f%u%$7IT~yq zHEm-GTW@VsGu~L*#Iv7v`}ET$-Z-%F)D}e@U)!8_4s)FV*H62#C()=d2tPG`e-xGjSbDc zr+!;Qcb}`@*U;Tt>JK;c3AO%YLwBw8|87G+vDSao(5;jHKWXUB*EzjW>(0~jUF|ot zUjNj*2HXD~+FVmp7uWE;+HVA{{Qa7{_ZBC@2ef6)^w;ML_ztK2JBP1DsQon4v}?a8 zIe35OI(VoXyMxwP+SabQ_ik$1yJ~D7|yX97FNIXIHfENHrf&js;p@%u@}_xSjqX>t^Ixw^I&z+<}w?B12Tmu}aWSL{#2y&sMJwYGhI&tdP^oyVf= zD(|($<)Qt*Im#IS4m-a0c8=>h{ehM;v^4ju-1|GV%;i0}>+E&*uV7=-wEBMpU};;r z{pH>dsu};|8rvuFI^D21#NYcx{k0FLT^6pt{g(%ui}i?o3$(10{@cPAppAVe_<{xB z9qu`_F!9ENjcpF{3AAeN9mksp=D(J1esZ5XsA->4V`JF=4!C>Dc-rr*?dtZI`^qbm3wd3eQ$KR z|y{}eF9KZWEPPpGym)!5DOYZm6CHH&hlKWk-+<92PEBL>TqX+rj@b1EX zAAGTbU&hHkH@utvFM~bzy{B+r{sLl5^Vj~%+U~tZ#+#Ol!%DS%1+e{(ukD`0$9nOo zj(%i$BDv4~T;aM}dymz)S{UCs_`;?78w%5E)yR1MoPn)1v-Npt$?aTnpNE#*=b$C` zIcLdz##wToZod)g`#iJcKFch* z&oE2AL4*4Yv$Sv2;3FE`XO*SD&nQd2O@r^);68_pzwZ@3e+;+v`D4j_-dOT;3vPXV z-dNhNZSWfkuD{P4OMjm?mfUBGCHI+PxcT}_G2GT?iY52?VacauvzFXvh$Z(qV#yaR zxcT~AG4>VVK2r?$p2Fve;o5zk7;fwH#Bk%SQE>fzo*28W&kjrO^TU$+9I)g*{|omX z#^-N8KKCoR&;G(~A8&A<-NkOc zKBEh__4!=MeLffNe(?ERxUJ9U!j0$ixo}&b!h2eFj%@pTU*f=WiwV`CG|-?pAW2xrJLF zpSgwG`phld`AWjAkI&q~t&h*#!fk!#R&t-UmE31+C10Y!ea05MpTm647Vh{yV+;5F z$Y*TfwmxGE*WYJs;rjcGE!_G0j4j;pe8v`T>oc}++dtRaP##UTe$i9j4j;H+dtRaP##UTe$i9j4j;n3Vg;EZakl{g&WUjY~lL*j4fP$pRtAO z?=!Y=Tc5Fo>+dtRaQ%J87H+;iV+%Lm?t-iPj4gIspRt8|Pw6wZaL4l*Te#!-j4k{` zxX;+a_4gTDxc)w43%B(dTe$u{V++^cXKdlxea059-Dhm!wmxGE*WYJs;rjcGE!_I} zj4j;y_>3*w)@N+t#`76lxbb|(7Je4oXKdm6`;0AIf1k0*H{_n&?@!+4e%jBcelO$q z8|!jI%_Xb)6PmW8xf_?i4_5QD>#?==0a)F8`4emHL%6zr>v2c#{kvM?{spYIA^v`r z{X2)BZ;cV3Pip=6I_5v%_ES&4Pr>%{`*1y=KOOzJQUtX|e(I=H%ij@SBJ0#5!jg8iIc=06i$U4K8f$m2f?xH;ch z(bQ9?*}&Fa-57q>QA>i7!0^nwh1<};Aj}`)}C5E5F)Uqy%fXyv@QLsK)r^UefsArA4!CB*SuP=_SEj3sI zY=8B{UlMFR6MreNe(LdA8th!-^Et47>ekZFj%xA$JXkH`E(ccgdopug9=rlaa`t;1 z`}*C@inVR~y5?VmtNF6dp5rUQ{XFV^c6_;hVn3HUp4iW%j=3hi{2t{?@GsEXS3m9g z`B_(6;(r;eR-P}b!PU1T&#Q>BI#|vAe%6(z|C(U6vj19e_0(`}u$uYhd|L;u?sta9 z^)s`Y*w4$JQ(`|W2d`h?4GQe%0-#TtD#$ z+KlICZSjWG-D~ld@T~O~;Ib}T!L>hLulLqqHTSV`HwUX_uWtuVtZl*BbK8TpCC(0D zwZz#5td=-CgA->baN_I&)|NQCg4GgdN3dGvwg)(Ib_bVp+Y_!mvAzOUODw-rP_w?) z>8oI$1Nhy2@)?PyeltG2#Z`NQ)lzG}hfw>N@0WUL{hrD8uT^I?nsy9_FYUdx-S+R; zeFiWJte^eF)_g3j@$IuuZL>9IAMJh|iMfAmw@uFdU}M}~=R6Lq=69!#E8myHSdJ@B z-$`KmhEE1h&*8dFp`FTc0Ed3s%vUY>nwv3g6XQUzF|)T10;^f)%vY_Puer$&rZt|e z^FNgKFpkXsklJpW@eT)j&V?TV-hd-(cqCXK_4phGCj0-8#n85I#?g|^f?1u&i_ofe(J7E#^)uTa~WJe_4NBXSlwrF@w*&uUfNvGb7|G`o_i&@oS)$r)vE}+D=j*_Zr|nu=`I?-w-eX@6UYgea+OpTw)Bi@W{WF)Fz-r!mCjVRD z|0@4m(X}Q2Z-CW&CI8#tH*+L^xqjyIO6c(2;_0IfXtFb{%f!^i&i)9xG}qJ5a-E)H`u zu3Uf5m+W_2ZI96An)xW$bHUbY+~c%QaQL$CleOJ8b9oAEj9eR@2CL;7FMo`~SjHBo z?=xU)m%h(})jBymU%yN1%kx#+cR19HBTk&}gUdB|9_|`gk14c2;CPNhKj$ph$2ldQ z?Hi2adF1f|Jn>%ymwCPfS4*Bh0w+&xKjcs|PjTYB3@-D01#a!0Bu~$eA8Ti9_1nc-oMRvz2?;EE%)`v9!w-1XO{=J?0cI%nJT|1)^~ zg8v0<4HECKV13k$_aUuX`hNsg3;!E<1CH$XkHPw=$LH_hvd<@Q=auWmKfwB^=bq?O zu(7lm$8+hQ;JrBF^BGt_bn$8P0wz_qmpVpUk);1r9 znsXE<&H`Y^Pu=C#GQPHQ>ug`!^j#2aJu>%&z-qZhXj8L)*3LSo|H5GFdux3iUIeV> zdOL@1aOR+GF^;7<%tM?wi-V1m8ZH4=D{HtU+!|_2zoo$H&fRCKa&vWD$8fCpeh%z) zC%*F1*FN!G25h|WWx=j_^7}koO+VvlQ;Yu>z|P0(&qLUj2j4}ju8-fZtN>QGzt4Z= z`um;D7r|=fJ@ra(b-z_7wryVjGvgi@ba3!HC!L{^w|bX*a5yRd>VLi*$%Fc`t{V`_k``i#?t1zlZ$OWPwoIVS8Jia zd78ia!Dxqan7^%c+!5XDMDp4R?!E}$8D6g0E^vL+NUoH$VOzWai$qjTM#*4HO|;!)p^V-t8E?Phg5H}#gk&RexD|Fzns>(;v1 z&#@n;&HJ{owb>_Qj{|RB%y~TAIp;Io1hAjy)KiCvVEbuvjDA|R#GC|HON}Oj{f^J? z4Yb8?3Rs&l4xp9idx-gRp5&oP)vkPtJ#c)%-cr%;Ru$r%2|Bi+q&S9R;N3Ngv7+S~j`s&XN9}C`yR)1~! zOr=#f7vDRM2df#+&uE$J3255lcVgjJo`)x)Ys>ZOYhe4TJKk}$YKeOaxIAx9h3k`R z_i5mhIrOu?T%Y7|1~~bad7O!^?Qm*2oVlC@R?E0&gJ)*{#eNQ4&AjzL7i=GO$32}^ zE#saK&bZ!ZB<=<1+7kCduv+XFfoCtqy%?^RaW4VeN8NGHqgAuNd0hrJmzlY;%D)a~ z_}2TcHrLIc0k`j^aP8@P1=x7$dnGu1wUvD@hilJ#uLftn<-WKEU0ddREm$q{y$)Q~ z`Fglo=6eI!KI+NqMsV`dUgmWbTzm4m8JxWQ+4YQj3%a)Cbt_mc_HTg8b+`?#mb|_R zwvT%9x*eRnw3m6^1lOMUcY?F8U8lR?=B{0SH`w*bc;5o+qi)W_@VN)9o^#<|@Esh+ zc75ggWX$`)<(Ln^9V79+4OUCM2f_Do*w1)!eG>0sa2f9rxcegU9tEo<-ecg0IP7OU zxjvcm6X3E2Pr}O@JO$TBJ!3o#R!@v)z}7naS+F^!j^6?6qn;Sw1*_*e{XMX8)y?y9 zTD9z{?}Ob_evWaCQ;+A-v}NocfIXiw&I@3*_`e7~jw9o~1lC7AF@6X>og*=R1Xj!Y z$ID>*sVB}WV0AykI&b&NkKyL0%{e_stCn0|17|Ju{|Q{5tldw+uX5;Tf4M%X$cjg+S_1t>zOg$ftyQn5BwfYn>G3!tz7Ð{!6?|I(^`#p%ZKZ51?UinX8 zzgM=u{j}@%I<2;>+52GEOxt^4xi&xleF)ZH+Xu9A-{+Y^^=I&}XzioT`@j8Zy$7>R z{Exxqp7=Z5b0K@;6Y$?S^m81!KH`7S=Kk(eu(56ZUgk5})~2L!>}l9=+HEuTv|wZ8 z-n#>?mhWZc|HN)A`-G`o|2qS% zFYooWP0yj`*y4;eqqbt~nc(GEGizro_3Xo0!1mLY+Rh5Lw%QVVHn4j5>|oa;*Md3V zYWleb+SIIZa<%<;b(#}fnfqMua=qtm z6Kip>F~XMsJ6`tZl3;z*-6QfvIE?LBV)Iyv*7HKVbd4SJbG4tXIW0^3d5+|?Ty3{a z4qpHpBRMP&R?9geUxve2<{(br6~X2b{zb5{9TrDwH z0hcxU5?r6`^HsrW?*FXmm*K9dw!~Qtte$gmb+ED3Jr~!c^<|y4t-+z@9L0&V7TEDq zce%BWudUoV+t)UI*9KdU%zYiOTK29swQ}vObNa6bw!WF?`rvh#r+R!g0IQpyKXWgS zZ6k2`nSWz#9LZ-Bus-U(C;2lgn}SC%hVQ}J6KgnFo8xRqE01k6aJgoi!_7P2xo!bg zbI)4)t--#my|%45cHppn;>6tsY@Dq9wqUh#?YD!w_S({Kd$78>Zb>V*#*XV4jul(p zUt+UOoE^dT&-gomovXV2N6@Ope;2U%X1#U=-$8Eb@!1WmZvUNW<+1GnF8l8ZzmNXv z@%ai^J$x^)d4zuzY~1{e<4CYR>RI=_!N$>M9=p@3CB`Ulna61OL*$_zpD|$d@Lq74 z$5^;|haU(o`y2#6 znLg_AIT)-SehAoF+#I#}U6Fr><0!DY zbvl~X`HPRC)n;8r;CC!|OIm%j=`)U2J$W4u-jL&&`aR$Tu$t%WEoiRgiE#D2haCr2 zGu}zG<+^_juHAJ%h1Qqru6{Cyn(Ho34yS?3@BL2K#*w<50rq=8_0;7|u>G`|$Embx ziFr0yE&k_#)x1y4wewtfz6;czzUP6p8TTw&d2APeckCSG7sB<)d+0@AHNRiX*cZd? zr_Gq>)2b!rrC_z>bs0FfGKj-NEV72==o}>;m)dT#W=d3nk+(WCD7~cjPH~c|x4@b_ihrnw7Tzq^UhVRai zef$VmKXvzw&ov$eA3$q-?dkg%*c{$xchOXj^MCr9qkBiYHBEhN<%`i;FWba=0=ytc z@^})g*2%ceRhydq7pArUqO{G=6#71e?)ZuQ4BXEPna{Iu=i{E!7QgR+wHf1STDjwW zm)7xYy@!8}_WK;Zv_D_lZH?ug`vKVZGWUu2dmQ#NjyPk!0CudL4=;k%%-{X?L$EL7 zXnTo6%{b!3`4QN8r+zPk)pB2}O)dVffYmbYkHKobFFEdO;Eb#7RSq@B6+7?DM?G)^8XKZ9F){j|mJ=U{EAy*##GfXf=a0ar_nehE&l+TP?)Ggq;>x?Xa1d1`2z zaef0f_na@k1*>_^#^)`txy)P*tA;b^+hCTe+MGjMzR&v|SXnTp0&%oBuG5$#_*XH>+4R3(CxAe_Wo_ot_!EfVZA8m=(0WR016YjmGdVIRT>fyt{ zWuNKb-dn22XL_)D_zd9rnSb`(jBtI_<1-U@ZhTVvnc@1Wr{64K_43|wR=7E8%f01n zV0G){?+`nGvA;j8E%%mlg1xuYN1HyMQX}=`H8;4tubT(%9!XB~!o9CkPww-9?WfI{ z{+_a0VlDtyi~oXPwer4hA$acVw5RXFU~R_r_nqajEebBg(eI;79akVn-(zL74 zdj9F3_+JJ)W@@<_IJMLkztzFoyjE^YtTn)D=DaGcTwg!;tp#3&I%r!HEZ24wwspX( z*0#05@{F+_c#YcT@1e{69^Pxg`rvnIjb}gY8E*rydaHi)muI{U!TPuI#3s*p8;NV1 z7{W{n}{E?_n7?u*^PsmHEhxiWUyK*XIia=0^1?roHGZ5<=Q-F4g=?$ITS33bqrTWWn0Sk1BhdjnqsJ9a)}$n{D8lfmWKr@+<9_wiHV8C!e$o(9&Iu}=r9 zW$ZJ+j_v!O^OozA{%3;AvCo35mGAXu!!x$_^gRcxEn}YxR?FDufgO7x&U?p}>tlcS z-38$6yYswX1T?J|z^yAqs!jw#o#JV$;1&%CYyXI@u><(b!Y;LPh%xmMi> z_F82;`)SX3H-Xj7&2i-!?`E+6+8kG&@ooVdN1NlybFI2n%;B}_HdH9obTlf#aWvmzAYJNsaJzs<;miF|039N0q zYP#0n^!_2(d7H1}Xm_vqcNl*J)^6=zrhSFO+UHs=*C(-m3@-C}6|UyreNA4k!IPKv z^!*7~Tli1GWvtiXYW~gG#QGULv9zb}&%xTZBky0(zQJL9b?4^akbJY&o!c+Lzv6Ii zc`uOblbF8-m-+k#uIAsYO+LScCm-$U`xaPR_}kzz*6-kI?R&uX^V)adiKRV#e-GA{ zx%~lbeD%!jk6`u8?Om{Q%X^JnpTztVxXkB0xY~)u+}?*LAMNS;0a#nEPalHS9NYKz zKZ70HpNIB3AlE1T{{k+@{wrL~zwMj4egx0h+SB)MU~L)uW3XDr{yW&Q{T*?~mg{4G zub2M-=X&`GSgy_MUVaLeYdaL19{`PO%+F}$d5@h2{4A33?5Ev)bnC!o%CU~Q@2Okg#~Se5>A zeXZ3j;MB^$(IwYrt!4wKR{p&%dB&InocC@2c9%Tw+jD|_-!`88v}e4z!0P7WxbloQ zH(38xp4j9WZyvC5v^lOk@7wc&9n<&i`DuOmzO6nVhnnx(;fqlVlP9l5 zz^Q|OhfJ<*1mi3Q&U>tXlT5D7_t?e3*4v!B!ScMvE&=vE*0GFfUwty>lC_>%$hG+% zyEJ%-Vjp}Cp7&U7@mmJ0%^3bYGCRiVJeLO$b24`RRH{#^lJU`Y2XJ7dD;^dC&--(!UmpQh%lGS>gME2EY1@oL&GSiYK0eRg0&M*9dG3~I>WRG-Sk3XxGjVM5 ztafX#Ctusqkp-N9ZDjb}gY8E+4;y16;7Jmc*N)?b_B$}`?q zz{b($xbj>N_X0bnbsI_R%etw5l|#*Wi_86@p1ekZQ->a~Jb8@)rw*gRa-R$Ec3ky> zU!t{-w(Qrj;Oy7TVIMU0?DKuWYVG^o_Iqm|JpG)jT))KF4?KeTC&vD8wcN{UQ;UB; zcr&!{ad3U|JDTI+Y!&{U27X2ftle?#qpjRG?w`b&AjT(oOoXc?4{d5?y;ot6r~jn7 zf1c4zhO3qR)%5qeG!>lnnF5w;^SX2pIO}sDSe|hX0cW2c43=k~9tw7!8qa>(Gu~lf zb!*|c@{D&lSbuGfE6;dGfQ_Tgapl>kM}i&GeR?#lFZZeXQ5VbRz;bOP80Tx?ylGTl~%hYcs}awDQ=_0;~C5_nE{z8?2_?eQ+)~^FIeH*XBMrADsE02bMd= z1+>P=`_qNsSMj&Me%f>1UIaEq`JQnxntJy7C1ADkyY5Tj>F1bo{W7o1z$1v8{r+{h zTK2m(wQ}9R#Q7rknXmieO6-Yu6*$kYt_A!0gu8QfxzDGrruF&MHMI7Zr~h?g#>||q zhpQzYZEB8hJo_6b{cjZG6MhrCd~SU+TpxA!`>nLT-0#|M;ZQTK*qYhj_hHA3%{F6v zLyS*qcpF@;%vr6>+1$$-T1$C;2EtaJpKI~wINl&fzZ>x=f5+kuVwJy7eJ5P4lbp<7 zn_BX{tM(7S8}9Fx=V!jZ1=mMCKKFnbet_Q#_jk+F=RUYT>Nl{;G}Zn5pRu%A>u)Ra zvew!j;83&HV&^-Wm~wUbn3~(ldupydxUw+RQ=Q6C7&hEKZ!K!2S+n_|x$Di=3W;>!Y5Wo&_f-ze`Bp z@1Sc-PTvLFU)`EyeB0c=eGhDY)=7VJGym_?CVzdN<4`kyv158Y`vExDv**EbZC7D? z5uEGU3t)NfJ6-}`kB@z{Wez_C``(?OZT}Hm&3g>h|0J$Bxy=Fg-(>MMCr5Jg-@h~m|9#8+TPWTDnS5 zt1r~xixk{>Emm`L?g9_v@Oz5#J!Cq#W52^lTwyvA)T))`oFYM;bmZ=s%Q_p?Sf?&1GX<@K)bpF;_9^0Z|=c~=S%f;5*IKHR554*wc zMf>Zg-JINGi-XNMdq}Q-a$cggJE!ElBv{>ZW@%bqo-^u8aj3bE#Kv*mKgVI+jq|5^ zotLTgth23SxIdQ#o0of6Y&_$}zFc7!XK!261x<+8d>rm==Q)3k-LngD7;hn3a{L$%6hHe;KOU4`mF~yrg_*`p4>M8 zJBBvn$Wy-!!Pd_{+ETxbz}l?eCbYh+pZdleYUV60(HiM^T+Tyo4 zSeyAdFL`WRfYlsBoIUNH@o#1OT81NaUAD&7&d;do;^jEZZ+TkJKeerir@!ZU#@hx> zJ$zfRadQ3K4z7>7pE=B9dohPOij$*%-^su2;BIiWT{)8L?qK_AGgtH31FUYY;^gYzlJf6J`SRb;O+8nuv3ak~ zVXkY^=Dlz)VmOcRo`QRg8B=im#uj{Ea9_~D^8vH9um`Qf>{57_fpA8qjB*Vo7u+>Cu;BU~+~9{6 zT>m2quK!UDeoVpjKcV3IpH%RZ!Kc(b^_d9Y5WjMb4?t7T8czbNWsRMqT53K8td=#N z3RWxE_#pTOjN=+u1G(7w>z_3~1ne5?qfH-csGji-1DE3+4p*DL_Hj=g3C?)dP;M;i zr+>yf8ti!bXw%0vP|tYBg3IxahpQb|`#9c-;Ed-Q$c<(F_0M=;13R8R+VpV^)id6y zVCS5lZ$BNZb{a=~&H(F^H98Bdb|y!B&IUV|#5o77R>nEE_DP)c;A+On^ZE0^=HM8{ z(8v9A0qunx?x~z#a{ZITMPTRWHE&(!eKAeu{H` zj=<*6bo<(jBj@MlH8%DZ9G;(B(dPWToEV-L;a3;j^XIyP>vuzg-&An@Z!Ng~w-x+7 z`0X{%`E~`|yvpzFiAeE6=y<;nu?b&PgtIKKf@r z-UxO->Z45`Yoeande*>KHtcl!M)7z{_YpR~{?f{#E zduCI7?*yw^!@Fqj=CH2Wb8>ys|65>V=JV4%U^V;SOM4%O{j=xf`q8`5)yn<$1l&5fp3X@wc0T$ipQpjr zrOf9UGT+Myyc<{H{>wENb$m$r|?y|X`U_Rc%REcecP1y}!|!9T2dYX5t3-807mvu5nx!C;lPdZ~pSs z_^;q@{EcZW?adlb!R~7!M{0aPjmObWs__Kc$sE>rDs5`~5i!dee_C+$&l;TWL;Pnb zxc)OW_{;_Oe4eY|`p?tg^A=qH#R{(f;tjq;%~RvQk*{+uYy2^qde;2!V708dHBif% ze*#v^n*Rf=R<8L!;nvx8vu1L!_0m82OhbOwz07A?GU zUR`jtod3Re41+sf&VS!KrbAQD`9D2aE$9D?;LOGKl$)Dt=@??y(SFHu7H~JQ_18~( z@|+c1<~bW&?KtLRUCeWKxOut;_A}2p(A1OXoM5%&IX5_Yx{u`M<{okkvHQn<$#Xt% zH?j5CPkZv5A6(|SAY85ev&HQ^7lxZ>p4*t`B53N#b5XEb^6UmDPxr6f+}yj4A$H%| zFL^ErHc$Qa)80JSkH+rnaE_eoN7VQr+9PXx80}FUp6kcZ=3HNjn4Uv&&#mPOuD*PO zuTXIPS1Y*wYc%+p1=oLrg6qFggKu1L{kJH%{#!Np)&qX1 zA60Pu_ba&m{S7{@;QCK3xc&z<_`wC&|LB72e{6#vS8)ALEx7)tH~1L^*Z;zT>wj^B zUsCg2tCohlhstZ!=g`!1?OX<|mTRYbN-gKqvS77bJ3kLrE3chjfV*!!58Q)tvHMT| z!E8)uIqza8g@S?YzR-To}Y5_^W1bCvFD}zlItd5bJbrz?a6ghaGC3HxY`07 zesAx!egrtVI)~)C8M?OQx;a=)ySZ)&POe@X>-O;E>Ku~m4q$D`bw{vTa@`r6T)i&I&ChF+!&@r z?hY<<-4m|12S;-K3fNpd7t8&+7rM6O`c<%6a@`x8T)igB&Clzh%QpPl4~DWExGOwPOkfb<>uE<>p0?ZwDwD`

t zzkb@2>jZF_>j7}Ji5$sw64+c9qAlw>8Lln4P64YW*8{=Hbt+hHeh1Myj@b7A`z6;y zz~-vIe%h1kq2Myt!{KU&aU|Cxz{$0&>yhZ%lIu}mwd8t?cE&ZQqrr0Vv9yjO_Wj3x z$@O@!x$3W<_T+j3xXkq=xY~&vx##;D*j(qDrrpo$)yZh;$@3JjTJk&%oIFnj%gya{ zTE`HdL2JL{c_!FA_18~(@;nP%=6McW?QD+Zc^=q2&!x4Wd7ckePo5Wm)sp8$;N*EB zSZ;0?(>jLu5?cEu&r8APslR^OoA-RDV)u0tNACH)R^#JnPpen{-bp`hx;_ia$e@}zoTl2hUe;sa4<$Lz!XzIC#xB{$}dkAw>%Nkt? zR?9uaRbaK8i@rBp1Gf&=)j7$<&PV^`b3NF)l=<9%rmny9m1j+E1b5?aOk-(JUN?cu zyl#f8AJ$pYf%RT!@!POsY@W*SOHGcr^80DIO8%;fX z@IkOz_Ml^`Wvw3qt7Q*93|1@m;G=NoWu44XE;c{?lg|@i=UL|SB$~SZ<}Xi;p8|K| zZ%kupZ`Swukj_cmvdO-D``{Xr-@nC__>0sKi}X#sCnk}4BRow zIX#P}o*I7#td<%(rdn$JU9ehe{5`N*S>x}+otJeoN4ePi^iMu7fSqTV&x>g4`kTKz zHGT=)jlVICrM+3>>#_T~h9fn;w#HY}UdLfwZ=lWf^%dfzuC}i>_-i#!-amvJr_B3D zXzHos%V0I@m~;2XaC32Pjwu&Ap8jITvX+_iPvG73(O*C9%{kwS-PesAne$CGzMA&t z8vA|!Ega7O8?>4K&xu*i|BVKJv*5nJyj5_YC;YD9j`2>-Q?H-G&7rK<>uBn!*U!Le zWxaj@cOKTo9OYv3(_d_EUT0IUU&6cTqrZOIlgF>X%{+dMrk?%$8?c(6FMKYz8?k;1 zSI;`W4NhLxLT)bBL0|p$px^Jo)?1r3ktfeTfcM5`fBR`~uIruHeci^9b^T_Iucp1d z#@o~0!Qr~zMVobfkC^4Ue$e0_7Tk6HxZplx{G{NH@sFCPo_~a!L%FW+qN!(H{{&Vm z*Y$n4^RO=FC>NWb{$g|Ub7$)HXLvV#^w&>&^7sq5na5wz)U&Q1fz`6ETNCSVaP_R~ z-@(btTFA}CI_Rt4cJ%ucY`wKv6M6FdCwLcZ_P1ZlUJd>{G*WdBhp+M)bvN8KzMA%1 z9IoZPv{}n(JBF@hXM=auJT>`@fs9kG4dE;Gtj#oVwQ_Ab;O64o98)fKJpILvwE%H4 z=V9<}`slBp_U2xG5WBDYII@@Tud#D}pvFE|`ZkC2e~32opOKj5UY@zZXDPUQd9H$w zhR;)Q$C$U~sn>LHb13UIy*?bN*9>5_vR*U6oriTXN4ePi^cS1k9!0%og?H0OfBm#4 zkJ-S@JZ48z&t9Gbtd_kzidb{P)w7r91}86TAvYK6ps#*o=r-so$Uk`I+T_35j>-uPoC(u5|;krISn{{27nB}@I+Te>7+;v^L;65u`rr?gT zY|T^81>ldOm+QJ9ntIlCA+TDxu8Y8(hjlSWx!C;l7n@sev98_lZu;o2pZ4UjIQUU) zWgbhQsb^i61gm9T`|)22uAX)M95{Jd3%R*i2YvM$PruKDt+zI7BG0-m2cCk>{`S+} zT-Rr@`+AZi>-toUUDu~;d>HLB9IoqkXtS;>60=;_l^T5Ig1fG(7u@F}YZlxw)~b2x z`31N+l=WO5O+D+n0$8nF*Du1IhjlSWx!C;l7n|F}VqI5(chg6I{j?{KFM*qRtcs?d zb^S6}E$ez9u~vhtXI<9-CogLuHy7)muYQNpZ*8#k)@DuQS=V*IM`N?U{j@jN^?B^R zzRQty{a%e-*XL?{BJKA%T-P7aW?eTRX1T5#HTcE_cU?Cx_;K(p3+@v}A)HiN5YUAF)yFKZz;FYBPMekagxYq0g!W=-T-*KNS3Vza;fv^UrFCG5Vg;mEqa zP-EBiMGn{UhqPJC9f(t|c7&UYb8}3&*zxoi zJJyNC+Ux@FrjP#mX>ZQ?RqVcg#F06_Tw~|_N{!E={V|90e~mWt--DRt{J+xRdllTh zJi6d#z!&?=j089H*c(kfd$|X!mc4vBu|~nwvzNz!lb5xSn~QbOSHCmqw-4BQYqKWu?B#vI z=VP7&1X+LOm5 za5Im|XzE$lDPXm%>xIOc3Rlm%9t2Ka)lymOk?kMFd;ryY*Z5oqf9oya4>YCaFicezKw)vb$T`u&aA@w8dL z>4|k5*gUm4w%qzTzvIEP(c0gB+Wom;_5570G2Z3y^$thQjo;VUbK?&+cHV#FuqJ<^ zwFWbh*BQhsYjRe>-Rl<>+%>tZ;Kska;QC)#^X$o!;O0=C8(%|Hw+_RwoeVa&{0{Xg zaDCK2qBf_|`to`-YOe10|JLs|_zCHYvg6sdYg6sc_g6scg&0VW);opRt_kGmKeS8O8?RF0L z@?EsP+{@bT z+o8PG&oSltl{K`Mxi|WM;6$i8 diff --git a/piet-gpu/shader/gen/kernel4.dxil b/piet-gpu/shader/gen/kernel4.dxil index c4446c5619a83ed5e5d47e218512dd827d013934..c48d59da9624714bca7fe6b0ffc02834646dbed5 100644 GIT binary patch delta 4522 zcmZXY4OkOby2q1DG66D#1cgo@nDE&M5hkEODaM3g&_XSCQKP%M13^Hg8ZB$Bjhin- z(5wkUZTt?lwdvZ{L|nyN+9o7ApjdImHr={3uaG`!yEiv!EUAX$6{Y&Y(vHL3Op*RTh04s|u=nT}CQT zT9yJ8@ z5K0=B;+gp#g0fVcQ90>w9Wb#iYNNuZ^iX)I@j=w#tUrin7#*S@`Ul#SgMYvj*{*JE z9aq-am3{}8F`<;A{TiZKWHv>w=bw0G_gZ&#A08psI}8&{ez9c|N2Z+L8BDWnD{*L5 zrD4?$oQ%{+nS;{5(8WP|YV0)H-cPiaio2<))pUsWG8*){oJ41t_*C(v= zjwvs*POR3qV-45Gm8=s3gjL*OihdjBB~BsY{CgUXy3y-Xx&8h@u42J7%CR3-H?p8Ndj3}2Bp?68&{IBz{3ae3(_z`f~wZQiYPFJgotmSJ^ZvMAhHj(IH5~8P% ziEDu)iGxocquj!8w`|7fsW2Vvy=wgER~KE}!C0e5Rm!!k@n6dgl6h1&`z_~2SHV>u zeHTk{dQ_@5GAWP31CPF${&1s9sN*3`XbCUaWT)dvewV5_$r(~N_h~|MHGdP3v!bjl z?~RL4?_3c&Jv8Pi|UP6cOxI+QMzeO!dyVm3hVWm)j+vh>fiLm zD|%>wL+ORPbxavX`!jQi^qk>!biZWxR>qgB;T4kIcwt}M<$1w6PE~RoC97aXtBGcp zoM0Ouk`ybv_t}qM&%>G+V|nG)C@87L9CbG7D5a5Z!?>eScS8fGI-*=E>O9-8!4xz` zY8LhKW;=_=Xw-?r%RH6z-Nj$y58^!l;IzgUN|uSwGg4mUeZRXMpz}Wf6$w+E-^^7v z`e;F~li_#hUhpX|=kycC-^#{NshaZ|-43p&2O}a~|BT_Bcrr z7IN_Wn=ij*-(h)aT2VZ7R4tsdahW)O7359I8CSX&_7g4dfT${>^AfiM zSN*_X0{Ac#GwJB!0dC2$i&g<{`Q+^y04_PUdKCayz5sxG3xHc?__Xx)0I{;2lks>~5FI6{U<4Qv^FAg~wn{qZ2eu;LzrXId1(4t?5Z= zDvuFAt+)d=e{)2}CA~vy`orQh9%Ih5~ZDW z0FNXGoNgD*?S#$K^*4LN{w1PDr zUudJ4^O1E}aG$eoaMj8<%RqAP}k(Y?@>M^JMnBtd;y5N_U!Z! zgCX&GAWoT9ywg*ru2=}f)}M(defa77rxKiWV(oaxc71!~8V)Kre zB;t7E9CiZ%0c0mr z0*+ac<4V^qkkGXtp^t$q!TVQedZS=^U$pM?5Q;SL`o~~;|CEU3skyN{Lp$(R+xpB~ zXnjjURJXN-GSAk9Gd2otQ+VV0gs2`Z+-q&=%l5f*Y+XJ^kj306(7GjuTrGVT$;O`hYF| zm9Isg)6>ObZU(lkE{c|qR{agN)cAS-eS_D49<9b(Kr)}{9!HPXNRW~Qc&57qB%O6g z(1JS&7y`M60aASBhf3#O7oWNrJy-1v>A*JvISMECayf zXZ{G_>1{g#xHS}6Wf(V*u{8y3>>X8%ZuH8hAp5UE07=LC3B64GP1fhJJVq@j321L? zQl6r8+pTIo?rQ@mopTofly;td5uoH`K4f`&n#dI!s?MdOd`G8*qVqIt<@KjC&*V^hphH za*`ASJ{70W?FF9-k9#3ip7Rs%@0*Z+XF&?laEQxu@Li-;BhiwFK*uPgTR6*00eq)7N*dOZA?t6-|D#pzooy_Xf6kNqankydw z3xCL>bDvt+1L}_}=Rgl@0Q7wDTm$fLi<-bsmK0fRc4kJV9pd#E+M@1j;Q)S;QtX2k zT5MQM5Rb{o90vI7PEv}@vLY+6edJ_5@-YRO=b*h&0WGk#+9WZ8%s~xQ4K5+$E@14J z71@CCStm0yv&XPE>b^53Ujp)Np{~tPhvR#8J2?0nzk@H@_+2UPC1{oco41 zSl^dQ!?~~eBiDYo4C9K9)=NOF^4YPvbh%pi(wH(bmX{bi?X`kBXtDUR039wW_L4ocJHDcIOv6I-D+_Kzu_a2mk7 z%qF&8ZA8bFhhGR~C^UJ9Xw`}HTgmE16s(a1_8&?|+jL{j$`=)^k5~CSn)0%fNF$}> zkrMKVTz6}w?pBrVmZI8Phz&=^m)?zQ?D#4#TZ0uW$1YT27nYO$w1Tt%{I5g>OyqZZ zTKi7gd)sKUHqp#Cp^mH2K{s^cbLcPK(0g;~=acCdhv?pZ`Z_;-*92Xb*3VBHJi}~I zuxOn464$qKY0K^v)@o_oJ!4#>{Y21djQmkX}rWB9n^@L7*Nmh zg0lrrHdw(C?^77Zu1aFl*@_#ofMaaWm~t_CiGVpkk_RLWOT_s`wz^S}3hGX|-@%Jk z3B>WXcC~OQIXDn zhIdu<()1;{OCs|nk#EH>qa3Go*d@R@Xzl<(T>wh;jUeiVC3D1bbG!Nze(JPED*w6H z`zYWHd#K}Y3(_=Z+5t(e9;vCZdT>I^{EaS+b%<}Up!xiDZK$>y>U+0^3>s^eh8p^J z^<;a~&hJRDJDPkYw563!Y2_2n)Fx8g)Fx~ACxu&${Ildg!sG#X=2rzRQic5XRqI8* zut{*sT4%t!WFN!%i4lyXR~GcF`VP0c(D|qVRtlRWxA1y{fugKB+3i}P>ehWjIZs~c z>~`HUe?;!oend8+g`v}L=XqBz+?x;~y$0_)z@1F+%FdIKmPGyWXeWNqm9a6!sYSbo Y@=ufFWCOG}?CS-eOQOLKDbN1?7h=8~bN~PV delta 4560 zcma)AdstIfw$JnMItd9#OdxuKJmg}46R=2q!~_ti(oR|#MXgRk6nro?t&d^aQa~AB^RyuXw89Wf?~xh?akMrZDcZj=-A$bIyQ9L`R)ovd%v&ukI5f7 zXRp0}d#~*M`>nNi4!mxAT~u0KT#}RDv$Xe%KYbPIN@Cn!^uCc!qtR>*Q7wzrGWWg9 zJto#_*4*^ZpJQy=H$RVhaADiN@`j<(>K7{8@@Oz+)O6ME4qv{D@XhKk&?A!egF z$=4>JwAJL^4I;Ihe_jzX;~`F7JB1RA`FU@-*lSi{p0tCzUy97H(#niAvnBHKk%{=)|x%)W)5|`T#+C z(5^5Y0gwNL9}{CS1(5Z-ysr!5c+xs4U_%+hg3 z>V9i8xtDV5q;@M?Bf%1{DddIv10)ffw3V|f#*`|IxyJ37X9%dTn`A@Gx_J33_LU^( zRV8Q5U6PGAv7g7IHeewiW~uONOJ`td2@`_}md0r)@lfs4MlrFa((<~ z(>7T}B}8pZS)XJ3Xw!0l0d?$@sE-Z<`1qOcZ1j~kKMZs|%zy(@#9@&nX|N#8)jxia|T9N{4T zl^XBgYNF5{R+_9vN0US&?!o(iNp@0A6x|>Zif5mHZ4)gj4T+)|)w*HvmD75<7=DB+ z(p`A|waxg`?h)~o3zXO*6o>bo-1751WDC9oPw3~rZre!Bx#5c)I9o_*<$?0LK$L0N&Vl8qoa}mSS~pW^7Ql z@EN-NOX|oOK&Q#Xsj-&?L=SpG6-((@*$-C+Zss2`{s-%eO~-*iReDk$S&iP+H-BEK z_6b~d2YpkJxhOTIbe@Dg#9gc@VkEh8LbakQdXn*O^m5*L-{EztomJ76k*ZX=z{z;K z4CT^Tl&@>vv4Zo?&vKe>Ny%(f%4YTR&)POm{)AIJ z??pXTOh~)gKSq|B`h)hrHwVlpYY~39wF+dmcAq}3;|U_b3$I)8vzG7 z)u(0VIo-$_gH?1zf)(sg%4>fz5`)4~?KK`&U=Dega>pGf#D48>p7+!hbxf-{6EdfL zsl3IpggR3!Yuh#p@co*LRD~I)6UuNJk7e5V_`A=hwGL%X)l;L&dwy*Y-!krK zutOw@EdfXSJ)hQelX9EivROGlNU)Uel=4Eqn5Z9mKIH8VJ79io$wH6YyaQOustCOb zEH$`p0ZUmGhSgz~ehw_H0+zna`8Yv06I<2X+iE@CI08C8Z?bY;i^5WTT=~K7QavHn zG>5AD5`qp2CNmTDkPJ4tU2gB*X6w0WiP)A)?>1RKS^c8R?K_h>S|Gli*uI?T=PyWo z^XeREytkbAII)?K7WuXMsjj+? zq^(yR_ZdUpqi3aQHU@K8ZAx(pXaI!;ms z=s2iOM{^%AkFBfQ*S6avd7n>Aw*Ax$0C4oq&i zP*xP)0TXJPf<7+R=S6sx)erN18Xalj{dQ>icJob4ni4Pi4A)Vpku! zgS>||ii`OvAGH(u9<)~t=+b&m7(D`L^YjONtiTFrFQA+J+Hday`X7KMH~)fePq~l^ zQ(us9$(!Yb%l&nl&C0%#ip9;IXnT9OoT!&GCaUr2P!Q~(lp?g;wyOc+F3Eccam%Lr z!{EyjuAclyl-}hY52Nn`+M|#^#v}>Y-^25t-9=bH?Mo!we*U-A^y_nO(72yMO1Ql1 zt!PLIc}*qijG7ri2S3KIr3jCkF~?>--PGc8zmyKc{QxC}@#A1ZQ@@V2!@CQ9U84!s zW=3uTrrigl2X+zCUOjU{xW`jLpk)v)ZN{^34~W)h4Kk8r&l=3e9DBPijC4e~x7*_CVc>-UMhXo4GkkyQG#3|jgL4`Z2Ps_DtW~v8`cS$4 zh(>KzijatTk5=76mqo7^hRtM*l91v)t&xtUpdpJ;3i({Qb1#*Rn0RM3PBWDZu*rhZ zt5xl$B#~4U)bAy8kVW{DY#ucqIVYU*v`z$t-WnZbGLsN-K{Dq0dDS&8_x?2bK}pb& zy3fjKA?iV5mW1#!3sZ;OO{Ba~5pbj$t(*j?i5`u;55GooT)%o26?gP2AlF{kr0W@J zXF@{qs=G-lH}4_^FxG@H6=-Stfht1W<*6LnWneU zk>f8$WYac4!_D`xmgxv(6`AaVup&rEN2?IPOICM-LZ+Ag%&;{LN38sf+LDziFHa%pBtW z9hf;ba15BqoHtUU{%b;R(BTs>C#nN8#~qUW)m{1^MZ#DQHhc&cUmX1<3xA9m-|e`WVU zrH%iAC;Vi12_aRvAv{kDA(b#E+CKGzMHjYsTYd@n<5y-uz^eG~Az;1l`sr}Q3m{_2 z_jy5wVLC+YoHg!na{SsE1^Gb-EGbAM_lY#91yGT1d;}0v7?RZr4)(BvDtFxBnJ!cR)^sqOkq{mD9Sz z$~7RThimPv%9ozXDNu{Q`%F%^57Re1Q`4^3koQE-)O1B;s|~A(Ss7N7A%5bCn)Ls+ zoW{d43Xq)de)TLWPPh@4(<4p#|4~!z!GA|h`loWL4i}}33z^@bx_hKh$0oYeG5!8BHB$`#QXu($JV%f;Tp?L;(T3DWW3v?$zzTlUgz7jBU&nh zn#sqO*g{?#oQN063#H?-vSurv{6iLX-qufXgRJ5OQl}7lJ6Zz+_p5oMV zwGxonzwiW-znri*N42E_Z1+vYZYy6;5*D6c>ylu1E{BBtA$K$>FGC-}Iozl1hm?cl z8*WLdibz6T*Qsnly0*;i>^VlhiQI=f(`26A)=9P?#|!3m^dF<%L>`FRcp&3Dq&23K ziKSFV!QpcXjztdZbH#5Y``6>qk>otM!Tz3-A4$&26Uw~^J95M? zKZ?vv7DIwkAEO6v<{L^YTG!xt`8Ht}j@r~F{jZ*TZGi&brz?}TVB(fXg*bOP{WvkfBma>;=caoz{)y!(%$>rUxc?G%CoFk>wNHH~1 zMju{AAFiSgS8MD=_;8G%^g+_At#20O>TMe36m@kp`*$g9pHe6!Ln{@$Eeoy zv&)VAv)RfiO2<;(sE`+q*@$`{7VLOC=L)fiq5J^&vzz<9kG9bzD(lWd&iemmCsO)1rzu(F^o{l9B!(F)x1fhaMkrI~6 zXz^S%f*UP}ys1>@Kz2z+FqG^A1T&ZLMEEm#){b_bb*K(V7~sp&_G&X68gM zPe#;U!}<(OxUno2aSjcvV+zqZteGjWe-9uy*^IeW{C-->raJ@vjRr0yB){N3{O^S|@VHwOR! diff --git a/piet-gpu/shader/gen/kernel4.hlsl b/piet-gpu/shader/gen/kernel4.hlsl index 408dcfc..0a6c022 100644 --- a/piet-gpu/shader/gen/kernel4.hlsl +++ b/piet-gpu/shader/gen/kernel4.hlsl @@ -1250,7 +1250,7 @@ void comp_main() CmdRef param_41 = cmd_ref; CmdEndClip end_clip = Cmd_EndClip_read(param_40, param_41); clip_depth--; - if (clip_depth < 4u) + if (clip_depth >= 4u) { base_ix_1 = ((blend_offset >> uint(2)) + (((clip_depth - 4u) * 16u) * 16u)) + (8u * (gl_LocalInvocationID.x + (8u * gl_LocalInvocationID.y))); } diff --git a/piet-gpu/shader/gen/kernel4.msl b/piet-gpu/shader/gen/kernel4.msl index c12e307..4d7068d 100644 --- a/piet-gpu/shader/gen/kernel4.msl +++ b/piet-gpu/shader/gen/kernel4.msl @@ -1304,7 +1304,7 @@ kernel void main0(device Memory& v_297 [[buffer(0)]], const device ConfigBuf& _1 CmdRef param_41 = cmd_ref; CmdEndClip end_clip = Cmd_EndClip_read(param_40, param_41, v_297); clip_depth--; - if (clip_depth < 4u) + if (clip_depth >= 4u) { base_ix_1 = ((blend_offset >> uint(2)) + (((clip_depth - 4u) * 16u) * 16u)) + (8u * (gl_LocalInvocationID.x + (8u * gl_LocalInvocationID.y))); } diff --git a/piet-gpu/shader/gen/kernel4.spv b/piet-gpu/shader/gen/kernel4.spv index 28d4cae6c7fba8be792b18fbda6f9c76e9dc1142..c38894159a01685a4f3a553e68d4dac6c41b07f1 100644 GIT binary patch delta 16 Wcmdnc%CZ57LOK}NZ4T{-t^xo&H3t&_ delta 16 Wcmdnc%CZ57LOK{XY!2;+t^xo&K?fB8 diff --git a/piet-gpu/shader/gen/kernel4_gray.dxil b/piet-gpu/shader/gen/kernel4_gray.dxil index c91d35df922eeb4350c523d83f7f068166f9fd65..739016756ae6b9bd43c267881383eef4991c322d 100644 GIT binary patch delta 4472 zcmZu#eOwdQn$G70$OMuII3eJKPf;T#qI`&?B%#16DyY?HTYrHl_(3tsR&5YI(f=*$zCG_Sf3>+iY8T-M34-wYsu? zZoibrSZ=OZNwgD2GDRyusf8){2cHi1VL~qrOk_=>2+-qKPi!lA5mfJQeMc48eZ?oI> zvV7LOW>e^USXGxtMN2^wqy320<}}hJ1wwCoa$BgX3YHY8$h(9_k4lz=@)O3|kP!qE zIw{W;Y)CUIE`vQPIj`^IrTDk%HW8w0QjQ4DQ9~k*&lga0)z3(95x#LxhI-H39eWtj zy@Kcmw%B949_lrsgeDWG9&0G_q%Rb4w~4$~J#LDFvEp|;s;i0UvMFP2Vu!OKUQz-3 z+LMt`{j*{To z`UGrC|5}TwiTvpmCnPqTSmgqg@5qzbIqSR|FK5Pj6{e*(ka!h*h{C#u?rZuO%jAabvMR3`2B>CD1r+e_$TyxLOtr0iek4gPDW$h@a_Qy{zc3b zhx4Jpax=K_D=W92j_&9BiM+b_r}z@aPSPDR3KzP_CjncoNqqt61mp7n{Rk-xzrGrE zI7k1WuRF8#3qY^!1$3=$HDkB!4gu&R_$ffQ7a3)uSWZ|^fSo7~!|xwRbvR%A!QM2v z<~(3GIMlVN-qqggvpPGZeBMIpgkFS3+MNVn}?WY*4qkRT{N4zH95UJ$?m=`IINl zBC%=u|9bN7lR=97j6l*&`$RNQ7IH%HJu%I;O&ezCd}gTz3yZj)vf;u3-8uR_4cJ4M zOgREEeO6oa!}Q=W-)lN}p8l*_idBRb^lppycvR;FTSu*LBH3du?`E%+>~X8)eS{NQ zqcyR1R9qy_|+&*h-|&L1$|ROc@-uJN%7pE@2! z-2(&j8sB$|lSZv>vk70rKKz5hEf)Gi{jMc*My<*w)BEF?a;Xqu`dx~AVAo_iu|_j) zxK)_N^aJ(Op1AzF!+Cpn=CEp`FWScwivsuj5k!rK4IC`S7FguDZdL&t4lc@FWOvi`mqMyShSOuqVm>+o z;^CZtNvLZ#yk{7_KKpjB5Xk;G9@Cq7Esr_Oc5T)cVhx#;F@-H={AlVK?>L;cM6-$B z6^|xTG|&6FM!Pz{>Z<_Tt(yth)_6W(s{tE%x3duIpAXnd{-{;ZZ2F>ZCpez~`&@uM zL9VU#7(m2%bXl|v)W<=0n(`#kFASJ)z@vIriu$Z$As!X6qt9v7f=C$_#Y(3P2mHR! z#UwOp#}0EHV0!`k4*G(FUJI~~)nP&bbW|EbTMIHS`582 zz0$PKBUI)g+K!E8Q?xXMe|!8*~eQtu-;a3QQh*dQ5andL3N&Q6;x zsKdkT2qT02r{_?}iE!#r$GX(IrjUl7VHGVJP5j)in9{N)eTK&29vA%q*eh1IWCWAx zhbejtb|(czZ%?7_hTWb;LDBbIv7prK&~efGQmT)F8~?jHoKQPDdg_5}JE`-4Y&UbK zf{QhW|I9`{8(e@1)tAA|->%1mSsvAe&u4)4A(}d-Fi!=YKtM@EE?oF9^bk;zJm;ib z6nkXq&IIt@5p@K;m+k#?VpGIxN2O#d)pLidAqwWHIx7bv&K zo=ZEPHUbQK0|uQ5D09x3>%46tCQNk#7rORNh)*41=bZX{`c+NlH9>?I)P=4EpN4?d z=MT!2PnXW}S(S0XKf#Pqs}`*Mycr#u+!H$>?j0?bLKbU%E0k%PzCwGwFgdX>%Jt}XJE`sLpg7Dl9>6<|PgKL!m`2SOX zR5kQ250=?lh{e!68Q_h}f2f1J;idYWo1u?*FV*knnJZqZzqYdEz_|LwJfME6#51A( zzJTg&jEtT?J%>Uz2m|UL_+RRGt(s83k_=M*C%?kG({1e?*7gB7FWrAl=@^Hn{|4LA zz3qpFc(-6#I_bKtRv?Mivan(j-|r%2czI$)pdxL|*N?(beAac_9R1hu@0H<}J6A6a z29W$!svq8~oCvkXDAwKdR3k?ZDwsx;rI@3==?0w*bP+agZpG2(@m&s3k_s`QbXX-% zK>3kluC6{Ok}N5pJr&C=CJuF1l@|0NG?bqN%H&gf-kvfql|S;Qr@Gd200$xAw-?b0 z-tdVj`%IiwSaAk*CV17TSY|eHGBmGP^y=^3?m=i?Oyut!^!AvM^=gn%?C(I1 z%xffWJV#OPa5g7OS_M99e3OZ_7pqcuR1MrfCP3l`oJL%dk>|C>>r5;zsOz$FORp$T z$NIJG|30&1R~OvO^H=cl%abUHpZO8|vePJ(+&PSc{74x?;z8I+CCA_KfedBhPUl^IY(0gAz+vLm3YoDI&hl7{~C7O8Gng~wcRQsQ(VXBWfWtj1zd z3g@6B*5Rn-t~GNT7iftECBj?KOe>V;`zjx?IfQs|Kew;8euJ&R*RDx6*RP_qwkDX44h+lq=2#(2avKrqY^DaZsW%7;a6hujjUXqn!xFmbSTw zW<2ynS;ZO3Z*7aI<*mvT=qkJ|dO`0&dYg9*C6@Xh0)4B9+N!K9OSXwiLbpXP1yk>- zD~FO?t3>tqZ2Q4I&Du(Gio>1UXG<=cy*_k%v_i&O0j%B^ToEL@NXdh%vO;%67xz-K zn@y}^N$8~8d6JAy9BU|{P8ma_1*2A7gGUAN{?+T5VL-1HEsHXeGfrKnmv%>TImy5ZSs4Q5T>=JSD+#?TkUALvnsU0jzSVO7&I0jDknmH+?% delta 4543 zcmb7Ie_WF1+6NwZ5P49rum?dwF>^)!fuUhAKJQv2P#{1dN~Dd$(a@(BbLfkL=VClF#|0st}r)&cA#rJf%$5dz_6 z&jK2iN)+@gWiBEUD|!x3R4_?}anCE1I#>uN5R59<>1DR}{t{6iq!Gnn7Ggsas7Brrm3n#Ko*GKgp%)_POHTE2 z*l~kX{z}aiR(X2OPS%6TeBs?HtUUBSZtc$>7rmwZ{Q2|k#l4fd)>NerPSssG_EFN^ z$qhW#89^UfhME%8nzvSe{%1OrXst11@+TnUK~;Ysp%j}*#SZxx-cJl(5tM#Ku6nM} zA0W-gaulO(?+zQwfm@X!wV%Wnrbbb>r@md4WBphX4O%H4Q_B5_wjXM@TSw7j5|GF<2{JsU-*!;HC>By6b5kFW5wUdi{of{tR!3 zPgB1g*~=a9TD(NP0Oig8$GX?W)E=@>!edcCPQ`nWcRf_E)}P+JJaVBRD(92WG;Ct> z51RB?{Kiq<*)Nxit3p&L-F%5UH~qHZPgGpYtD$}#Lw7^~nylmdWAH{&E;(oI7AF7v zGq`=gMbX^xK2R=s-$WpR6hKLz0GiIP^U+52@%aIT5>YWp{~~C#T+<^p)=aDp8Ih-N8Es}k5J4_hHvDpTu*UC+?I`uSZ_eAOL`-0kO+P&t{We{cQM(XJ zO_omKZ=>#MvoU6XNE(u`!XOXV^is!z+T>C9jf}9m^jmKL-UmXR0~~&nV})ZTTr*4s zywx{e1iU3zK3oBKtJjf(G2$*_e8(uC35a`DHI?}GsC$=X&Uo;}{(ZPcJ0}{9Y>c>x zcSCAgl|PWuZm~@kvH0A_i2*Lfj%}>Qt#=$z<#ox31Xz=>gG90+fRP%g@!Mr!JuM$3bmv8pp$Mr6 zHEjXgD=h-s%OA=bk)+bV+33Lb=GbkQEmnU(l?|d5WIc&iw%cIt6FEC%^wqgRQd;71 z5A{pEnKa!=%y)LU=+OwUM@d*Cl$ID0E9$yB#PqO?-jEwK>aam07B;3O_Bzd^{4U_! zkoyn}PstihkDXn>`Js%i0nVa95;yTn4_G(?<-34y5PXKfR3pR%3$Jv=lKRym_#^aT zlWr`)f$PwnHE}b0Hhc(Wn)GIG%Ot-yffT>w7TUd1oFM1I8nheW z9@U7z9QX?q>5FHl2pU|eAjk8&U#-|Cm&Yqdrp+)7t`PGimKqs2`J1vN!{l;YGucZo7Z(Gk?IS1QeC;k2CXEF|T9FrcQ7^ zcU(zcZR7six@JI4mv0`qaQclDKK*sjl2CM6e0tWnHu9f{QMN)xVJPN8EZI2*H#0n9SaJNX5}}pXB~1GR^grih16= zPqVYLBNADET=#?flrUbETOnG27TJH-iI+r5J1cY&jE+uVh%q6$XywB zOjj@sm`cFSDYu?tWV_{^q+)xQJt;e{8){Cx?V;Y4b(YL=%IDg<2+2`da(3PTgayIW zeIO@Cmp&?lnZ=CE~q6@ zcUKPHV_v5-*s)09fj?%yCu37-k;0)Bp&AT-4n;fP|)h3Ub zpUUZvYC6sduM#JC7sADyFn~Q$Q^h*?Am^!?PI2tIx~Fn_SSQ}|R8GYG#5IrPNx)&IZ;4O>WWGu7YOyQpQ)2c zqvt*OCYI@u`8R3wf~S77xxPe7xDd*RCK#hiLSaNg)P&Sf&{OYM5kYH_$71wLQC*{F zhUsrerSwVddI9EjAmS;2&;WreD%&K%DNOtFkt1#n44&{F<7|94(j-&WQc9{CTn!A z_?Kr_A8YYiIHs{=Aq!o9I*=g3%4Y7pF~p(xAVU7GAAj z#&dIUjq{+Fg;XFvbC{saS2mWl>8DLENN6@B`UtebXw$!6ZzRy3hi>d$#cXa(u=WnA z$*Fwde??JUTWH$UrOqPp68CMf#c=cCAuDnWLs~ zZ%+X|=bvzm>p(Q9FCBGfHi7my>1+W0-C$!pg!tSj6}w;YZrZT6RC|Po4iwnS=96C< zQ{dgvj;NAVKH`>}0YeagNZVm!sL6cpxdg260~S6|0r6UFl!k0R{mlfdM@N=Q1`3M` zypFOrQ#LRsOnA{Qf3YZ!2v3!}%eY7VX}$h?DK&%%J>&vOm0&{ZAtLv5eY-5Dj#j%8 zT&k;0g)N||s17>dc@qyTBkdL!Fz{T^r4-bxO7koN?;5pz12ZfKqru&9%_9fLWIN60 zuf)vnwc>~{UpVeSuC?OT=DU?xPU)s>z7z)Um7ai=h#f1($p_-7G8Jv65_l~!W}9}Zpq;Av8`*=)#58)i5-YUtz^fIYe{#a3 zk`hK(P3D{pBo^$jacV(;8P6ij1?AV#Y)CPsL&bK=4q2@cS=m-rn62}$D@m=i0!woZ zG_=b?L<3iY@mD&C6pe1pP>>DGMyR%+=Ebf|(BpQpjDwcq2_l zTc75cYB4~d<*8P4U27Fb49s4QmKSh;W|)j|J*wW!-fIM;j-%= 4u) { base_ix_1 = ((blend_offset >> uint(2)) + (((clip_depth - 4u) * 16u) * 16u)) + (8u * (gl_LocalInvocationID.x + (8u * gl_LocalInvocationID.y))); } diff --git a/piet-gpu/shader/gen/kernel4_gray.msl b/piet-gpu/shader/gen/kernel4_gray.msl index 04f3d69..a190ab6 100644 --- a/piet-gpu/shader/gen/kernel4_gray.msl +++ b/piet-gpu/shader/gen/kernel4_gray.msl @@ -1304,7 +1304,7 @@ kernel void main0(device Memory& v_297 [[buffer(0)]], const device ConfigBuf& _1 CmdRef param_41 = cmd_ref; CmdEndClip end_clip = Cmd_EndClip_read(param_40, param_41, v_297); clip_depth--; - if (clip_depth < 4u) + if (clip_depth >= 4u) { base_ix_1 = ((blend_offset >> uint(2)) + (((clip_depth - 4u) * 16u) * 16u)) + (8u * (gl_LocalInvocationID.x + (8u * gl_LocalInvocationID.y))); } diff --git a/piet-gpu/shader/gen/kernel4_gray.spv b/piet-gpu/shader/gen/kernel4_gray.spv index 93ce73ec179103b7d6f61e13f1b5e2ec6d638a53..17c75317fdc52c79bd1ffd7073f75e80d3b670aa 100644 GIT binary patch delta 16 Xcmdnf%(ADMWrIKm= BLEND_STACK_SPLIT) { base_ix = (blend_offset >> 2) + (clip_depth - BLEND_STACK_SPLIT) * TILE_HEIGHT_PX * TILE_WIDTH_PX + CHUNK * (gl_LocalInvocationID.x + CHUNK_DX * gl_LocalInvocationID.y); }