From 05dc88b70ffd7fd8d4437b487014cd8f540c9013 Mon Sep 17 00:00:00 2001 From: Raph Levien Date: Wed, 30 Mar 2022 07:27:29 -0700 Subject: [PATCH] Fix is_clip nit Not necessary to test is_clip when is_blend is set. Also recompiles shaders on Windows machine. --- piet-gpu/shader/coarse.comp | 2 +- piet-gpu/shader/gen/coarse.dxil | Bin 10568 -> 10544 bytes piet-gpu/shader/gen/coarse.hlsl | 75 ++++++++++++++------------------ piet-gpu/shader/gen/coarse.msl | 57 ++++++++++-------------- piet-gpu/shader/gen/coarse.spv | Bin 52356 -> 52244 bytes 5 files changed, 58 insertions(+), 76 deletions(-) diff --git a/piet-gpu/shader/coarse.comp b/piet-gpu/shader/coarse.comp index aec2936..454371c 100644 --- a/piet-gpu/shader/coarse.comp +++ b/piet-gpu/shader/coarse.comp @@ -305,7 +305,7 @@ void main() { is_blend = (blend != BlendComp_default); } include_tile = tile.tile.offset != 0 || (tile.backdrop == 0) == is_clip - || (is_clip && is_blend); + || is_blend; } if (include_tile) { uint el_slice = el_ix / 32; diff --git a/piet-gpu/shader/gen/coarse.dxil b/piet-gpu/shader/gen/coarse.dxil index 5770e6fcc85280788366ecbb78c09d791310e625..12e88dde300ffb33e17ea6c31c7c14b60812b006 100644 GIT binary patch delta 3829 zcmZ`+YgAKb);B5rblgs9@2F7AcE5G%8-kvCRzv zqLqlqFjP$xkZ_4c35tyNO8^~II!25wDt3aU4l>AO5EYO*b5fG7Z+&Zh`SV`(dEfoK zd++DjE4@;)G-P>}C|SRMSFGSa!b_InYpXVY_hK;$008dA5F!TaggVbK+oCX3EU!!S zf$@{iPV#}Ci*=#93NA0tTD5v00sv$H7I=*V0QVX|PLhzrR%Bn4Z2d8k001v=$74|E z+USrG#K9l{P%tms?vlDuYlideM4uq#7Xa%XG6EsMl@Lp}085{5{XO`CD?#aS@^6I8 z2${e$)X7x$IN}*cQM1YJ^_XY0pe9q>{!ab_*3Ut&SpaBM5dC9&+o$kC6pDNNa7c@S z)Zh&q$7lK0+{a?C^#s zn2kF~2^)oBf}mZD1sJo|Ry6SeyqJrw1CGe50sr#(W1u!v{YSbc8|0>HWYq%y6aoRj z*aVBhFD?2eTe~D5mlY{RDQ-rgLJ_>OcJPUN#k>UnX!qo<#i+3(A7HF<&6{ONj7hlXq*j?S!j*x<9RlA5MrhHp zOjn##*eAeV+#LiYxDWgoTBQmrR!8)x*bY^6uPruHm$Xs%*6S^Dl7CP^>2bTuk&Vl{nPupZzPLeK z7={7*As_{(bTi7M_(PIk92IYmz+vu(SBXX_yw{6sNkEV;AXAdG>Bg*8o=h*ZMfrp~ zgxqarfQ|>N{BGQRer>JqzXSebv_}cU37J{^mALoeK3i0G_<&GiGABuF??%dmK1u~0 zmd8#ewNR9#StSXtBE4{M#sV_!NxW@#_fWkJ_(QrEnUo}MnjGg74t$2o zjl|_0_`QQp2DN`WYq$Sv*lIOcn ziw3P=0tVsB9}PHLnF7IpK;YGjE-!AC5KVT-WdB!{T(n6D+1&UL+eba34vNJdYLWQ0 zOQx!nBCZEoK)Yld6XO75hD9;&^V{o5edI~+A^xzabN4m7?{&N1b^8=zQ0@E+Lm#LC z0U_7<)kw)+oGLP09SJMckrPdt5(|2s+GFfuY7Z>bMe=l!A(fFYuc8KMIrtO;4UrCI zXkEz6DWw)eI~z$8>C3m^^L%Gy1|Rt(@YqKcAojP8q!37MjagPlgY-^q3I%PzhpM3+ zsVnC$SjPyD)5}Vw`K5dFOMh%6gsK&g+-v-lF<+!Nl$OsQJu|;tpRG=S3uUEy5CPL> zOlf32GL(QsU@omaGnrktWFj)6^Gu+|2YHQdn5lLAqtxw79jt~9rnWdE;X*uSQO$Os zM1nvMc0=!EgsUUMcGbnP>k?}h;7e4H-#dXVSffa)*$!hf@VDpp;@NaY_@aD(gz+I( z*vr0B$KO5^FRqIp)pz0I)KH%n=Y7VKOM2PYwM#bDEt%056d_Sa*=q1{F1PV#t$fi| zL-#~dhASzF{#QB`N>+V_DpA3O_7W)?3Bo#ViLA%q4#LvRE(V!4F;X?bte7yCh>Z-B z^QhOhzMnyXUK{7V&aZeqQX)1p;JGIio_IY{jOIVW@}DyJ>NDzb1h#Z*k7Dk|4Nnfz6jAGgFQyv$-CC3)=5ABE+4v}?q0%4Q5|7aT=jVYQS$dY%T9vHg ze2vJ`5)&rKQh5}PN2CHZFa4q;Srs*18Fi~N>dzS;{{<6tt#Sh?%|oBC#Cv2It7xBe zNTVM}X(s@$At&*H6vke53{*xxtfW2oZU1EA!%7%?#WL6*{iL53Bt$^{DiMQ;Px@hO zN0pNtGeizuzh2BubAj!|pNfJ8$sUN8)34^H*kTZfR_rrkf(84D4v`{GtB9+ZQ>@I7 z?p4w*QkMPFA3M}vWR;xxDuUW+LO61V}m>U@} z*N-Runey>>WJ{NzWI+XrNcdz$q7eIYy(c9miXM0=eL&heN^+h?)0&qbVKHMF}5fhhfoBZx5 zU&TvbWv+cyr4TRZIlh6r))UK(*y;Wa+57-6$1J9;%;QrJuZYV${6Q z1n50v={WI7(YOwpy0viXkbcUyx)8Vw9`x$qre?jA?<7>&RjZkVqd?YXmvy8I(A`hBx zcum=mq(Co$8YJa^VjohljD~uP`ncC?k+@+8E2}BDVM9T~S527T_V*Sv6tZMZnue_M z25VF8aRf~Q^YZ?~;f46U5Fsz55H(Nz^y|~^ZMbmC|L<wzycoyYGQrZ9 zH`Pd+JhFMay#Bzda|*X^qIGI$o492^%cGn2lFeS-l(tmU8c#r+lU=N~+64oMZr=IO zXS4G@eO_M30o6RUxZ)c%R4Df7`zI}@c#X=&tn$YF<&Dm!?|(XHP@nAgG?qurjr}ZF z)6K?%w;P9>YO4{fTVTZquyl(vJjZ!zGdyDMUF^h85ah!38e|0!wz$G`T&FhIYn`;PvW z)3uGA%^ZuU#Y|hY)GhLdF6=Eot>|~>h~2?50wxR8rXJD?NJ!}9q*8C_?C){DBy+AE zXt^@ll8xR88_lWlc-$U;>(**j*v&{LX(z{MHFcHd-%UJ_xau z98Xs{Y8ip*)tv>S&F4qi4b71VUb>=GMO233f@ZKmWMrC_qS>jt=bYzSLPAzn7#BQ) z4JManYj^GU+1jx+VYMHjyn|zOp0^z7S7&Q?b8*`|J>mPOoGS86LHkVeg&Foab}z>y zyWs5!KTL2=K=0Tey&JvouHc;5&hhkw@%DF$f(EpG#@~*Pwx4a1zJ~IA?OQ!J4xpo} z<3yYbB6rZvxeW$1sZW?#dddHLZ|&!CESuP{38{Ui$YMa>^X7WT9ZN2Isp(su=|0bN zik*zu`n@)}Uwb|~)%0Gwl+%7dQBdvl*Cqn9`!zfXE-|b0(y_;LMcfe>iqbp>2Rn}$ zcI8@R$Y*)Jj}Qpf?*#s?vT%?MJSI4RiUDNH{C~H4tN=p7s+VljeOjSFA5YT^fgkQ_ zT4WU1v%!6MwTHRAX_sXO*q8G^M@m2&x`FK`%|0cP`CmC~dD$Fot~5QSt-&jYz#6C* z`Cm03WJe8!&weOJQnFuHY6u%tL-I6F^?w@sqo~Dj#nEyFDVC=N)q?q4I@m^G2B1DD jP8WcmXa{_IWp_w?RLG8y?-$ZF(aCQyxyV1#-+uoE7SY4Z delta 3860 zcmZ{neO%I6*1&&1UKNlKG4UOiaY7V_vKP}*Kr;tTk=QA%W|1jZZO0{bs%f1NNHnwX zEu+R(c$5_1!VyQCu?Cdd#;j0j$8D^%X=lyMamprhtl1w1^Lakc=kxH-<=%7eJ?EZ# z&iCA(L2zB*9k-JkyYp`=`px50&$A0gcKzycYZwFoAfeM6k3c@ce60=b3Pq5Rk^d>k z-(6Rs{%zNugR#7&Zb{tE_mk!V0YC&Gfh9Bm*!uvqk^n6pgT}eBru$SJ06a%mhk+Nm zD41pxIsgR#9*7rRw+RE_w)xCCJlwPBuR!q6Pz!hiUwiiru#uDiQ!eT!79XrY8$LHj z;e^1CVBK|lJLPDn!KDFvf5y=o!TIZ54~U;4v11f<5TKA@G@ET*ACXsqAiMg6cdL}G z+eXH1NXUVsJdwLOeG+%`tu=KrZ^VF`^7~_)Kgw&)!pAhqnsX)}^NMAVoeGwQ%g z$IX&h)&>#lh6a^Q5Ylb9TKfEnf8Imv$YpKVVfJHt7SMta1||5(FYl7WS+EA6TvW?; z8w{93DShRI6h#_}wNoLgl@__h<8T1N%;xxaaXwE|u0M=UqizSUh01)PW9+?y+T#bc z;{T_mfA?0N1hUxqrS2}P_ z5^IfwEGMAz2vQsg1)$O6e#VN(KGnu!x89o>V$SnIhrVUZ#0T511~$e8NS>Iw;s3z5 z!tuMqoYO6;NQ-J?`K|Ys2ANB|(D`eO`S{>AOJLjY(;NQx^tMxxZKpO$zA@c{d11#H z5}oNacm_QlIWFZ#SqM>$_=x!=$KK)%XmJx>{AaxQC02Y0EB3;Qdobc181YNA*uG?l zA0?Tg9GozOUt;mez)C+-KEF9O1Ap~BQV-w4?Zj*eCQ12^IGy+{QKT7uQEVj^UJQy3 zN?13|jri(~N#C&U^^>rpj3gcFwOLw&9*sM z+gwzGE67LuWGV$=MHF(@GkEmJ9;LUlxPa~7(K|TM%DH^cV=KzvDpkh(7`_2qgp{(~ zdz||`O<5e9$kq{lEycT5T3?T^SSly`*$-~yPKQS> zgr6e-p6V4sNlf~d*Ct6UrPLg%@E_-}#0G%kf+FE=rc7IIyS`HWX0V2u{c*ybV(FjJ zbh6DH8XGXd;kR33_?+LSN;z=o)RNUCZo5Zm#Zqa7%khe*wvhQ1!nv48vDLoSIPRqN z3E~}F1;MqFFe?m?VI@|e;r>7gdM^jP_W<;F<@h8xJ}h%X9b#EQ@?p`+0@2DtP`08E zP;gt5x?>Bss0l(5SDKD;GlvcYEP$)WO{k3sC}BP`q0J5xHd(m6;YKUIn~&}0FVs|U z)6E#CEcy}mD~~>_Z=V%=*}71p;j$d2t+;HRMX+or79y4v7v0@P7$Pn?7l{zDefbtX z<|-d|mH!G7NM#2u)9T$7D1gJ7=%*`>pev~XRaEaP>T~@-u@Od=J7m|glt))-sF50~ z&;L;8FM(q|BA0la4@8clkOm)2y-A7i{etQ%s4o+{WV*f)Y2cCt;0fHkg6C$LipLS` zh)Hbl^{Jk|WI}51Vk1umBP+%kpV~g8GS) zvcN@c;JNxVc?3CERPqS~Fr~nySY#1_YyytHu|@kLyJ0b1@uxS3Pvj74Do~*K6KPq2$SGz8~VclYVKRR5Fu{!PifwsQCUUTx?`kaRK z7;VmB$P2s-4SXWO9L+XI$7g5;h~^?v91B1e;GiCyl>o-DL1|2UMy;9jkloB5YiOM~ zDY5|e0~u4xXurAsn6;J)gs~1q{DDk;*m~@g1{OJz9_O0roy_qpokZ=42x6)CB_M^) zi`X8td5tzhCCFC&lBN2v`B>Bl=9yFWQBX8jt5F@%sur77;%3CAw-+mAO{D8p0{UFT zo;g~~GnzMpt`gOXJRMh?9~2|d{-BzjQoWf{WoAKwq%`WQa`IXz>A;sfsgdhK*AwY7 zB7KxlC2DP48U46|e!PnQFRPpBpA%c10UCO81dD8Zr-B$=Hy>Yskg$xc z7bA>GzLjUkO+cX8fvt&0yeN`s1mcMBIgf7+dF;rlAJ!(ukgUXrVIp%_O7yuin#kvb zW+S_x&Dz;uJs73Osz;v z4EwhrtUJ}+Zq)KVy6s`7Q#Fal!yZ?KT}u696QT=ofd7P_yn_?wH1Vm_n3zSHx7%oE zUFQRaME+|u)Eok_?iX?|?60VHsP1A71lyiE51}MEQtE(4bo;13 za>kN4V@aMt@J@UE2jSURed_H!f0y#KW~(e}-Rz^}*>`w%=7~c@OFtWJGQ+MIU|0J3 z*(F=5z|e}lCd{-q)I*C*@+$^8BnVZ6T4XH0)vw)u_5@S-;HGTvH;67=mk4W-D7KZ8 zpvQvL24gz?b+%^n@exT%&?hM`3sPEz{hP9pFPt(hDY{Znk+QKe1;&y6X+8dEN=g#_2mR~0%!dzb#R$e?4= z`lPb@h2!-iJ-%=RGwbVePI?HuoD zN@tvbor&B%sN2n}Yk5Cd9aO3ipXeEzRBD<`qFSm~0hNA#VR?8``C!g%je6fg@g67t z7}AiMXF>)CU`wLYsO>~=qHOPW5Dl>?!NK?rUZT`-oZ<&#P_0eUiiIQ2VmzdWMP z%h7j^1fS=&rpZ4fHyaAVCy;!%gBi7xUWL@=`Rl>Wl%@iU#*VYzFQ@edn8HC_;b2oC z@9#OA^DvCDU^`W%kiyyaqdDhch8z>7%~baHmQNvL7a)%oJxAUjM&5r6{l#vE|A~E@ z?X`Q|l-<>H$Ops72ak)>P?A!(R%wu9`kiiUZT+IG$<%gjux)6tZC+p2Bd=0@C5Ng; z^Q(nma;`zOXn*?V6G|MFLrTB0s6X;m>(_#iZl6|1=~!xMn#pcoE**7pjE&dI4snzY zGCI%cIv-}&!&P65wDy*@eseB}QX<1(w2l#sy3rBG=xzP6V|KdZgBG7Ic^WLAJO%XT@wE~@HIF=H z_{&;4+SjdL_9r;h1IqJ6T8W67r`{J)5qrf1}?<*>ykl&{6&Au$ifE!;zAt372>R?dPl%hX{ zZvO2v=HmCnkh_EJ!^Z$}ZiT)Gm+AT_L(@-u$SHsyR_{}zkGKpDGn?5xj_8kg`Hb0( zIeOyq(I-7e_jhEvPJwVSz=OU;cPYDIbQkd5nq)R=m-4Rbbo@B6DL3TEsQ LVlOo5_uqd3f3?N> diff --git a/piet-gpu/shader/gen/coarse.hlsl b/piet-gpu/shader/gen/coarse.hlsl index 57b400e..a702df5 100644 --- a/piet-gpu/shader/gen/coarse.hlsl +++ b/piet-gpu/shader/gen/coarse.hlsl @@ -852,23 +852,14 @@ void comp_main() { _1551 = _1542; } - bool _1558; - if (!_1551) - { - _1558 = is_clip && is_blend; - } - else - { - _1558 = _1551; - } - include_tile = _1558; + include_tile = _1551 || is_blend; } if (include_tile) { uint el_slice = el_ix / 32u; uint el_mask = 1u << (el_ix & 31u); - uint _1578; - InterlockedOr(sh_bitmaps[el_slice][(y * 16u) + x], el_mask, _1578); + uint _1573; + InterlockedOr(sh_bitmaps[el_slice][(y * 16u) + x], el_mask, _1573); } } GroupMemoryBarrierWithGroupSync(); @@ -897,9 +888,9 @@ void comp_main() { uint param_25 = element_ref_ix; bool param_26 = mem_ok; - TileRef _1655 = { sh_tile_base[element_ref_ix] + (((sh_tile_stride[element_ref_ix] * tile_y) + tile_x) * 8u) }; + TileRef _1650 = { 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 = _1655; + TileRef param_28 = _1650; Tile tile_1 = Tile_read(param_27, param_28); uint drawmonoid_base_2 = drawmonoid_start + (4u * element_ix_2); uint scene_offset_1 = _242.Load((drawmonoid_base_2 + 2u) * 4 + 8); @@ -914,11 +905,11 @@ void comp_main() Alloc param_29 = cmd_alloc; CmdRef param_30 = cmd_ref; uint param_31 = cmd_limit; - bool _1702 = alloc_cmd(param_29, param_30, param_31); + bool _1697 = alloc_cmd(param_29, param_30, param_31); cmd_alloc = param_29; cmd_ref = param_30; cmd_limit = param_31; - if (!_1702) + if (!_1697) { break; } @@ -929,10 +920,10 @@ void comp_main() write_fill(param_32, param_33, param_34, param_35); cmd_ref = param_33; uint rgba = _1222.Load(dd_1 * 4 + 0); - CmdColor _1725 = { rgba }; + CmdColor _1720 = { rgba }; Alloc param_36 = cmd_alloc; CmdRef param_37 = cmd_ref; - CmdColor param_38 = _1725; + CmdColor param_38 = _1720; Cmd_Color_write(param_36, param_37, param_38); cmd_ref.offset += 8u; break; @@ -942,11 +933,11 @@ void comp_main() Alloc param_39 = cmd_alloc; CmdRef param_40 = cmd_ref; uint param_41 = cmd_limit; - bool _1743 = alloc_cmd(param_39, param_40, param_41); + bool _1738 = alloc_cmd(param_39, param_40, param_41); cmd_alloc = param_39; cmd_ref = param_40; cmd_limit = param_41; - if (!_1743) + if (!_1738) { break; } @@ -974,11 +965,11 @@ void comp_main() Alloc param_49 = cmd_alloc; CmdRef param_50 = cmd_ref; uint param_51 = cmd_limit; - bool _1811 = alloc_cmd(param_49, param_50, param_51); + bool _1806 = alloc_cmd(param_49, param_50, param_51); cmd_alloc = param_49; cmd_ref = param_50; cmd_limit = param_51; - if (!_1811) + if (!_1806) { break; } @@ -991,27 +982,27 @@ void comp_main() uint index = _1222.Load(dd_1 * 4 + 0); uint raw1 = _1222.Load((dd_1 + 1u) * 4 + 0); int2 offset_1 = int2(int(raw1 << uint(16)) >> 16, int(raw1) >> 16); - CmdImage _1850 = { index, offset_1 }; + CmdImage _1845 = { index, offset_1 }; Alloc param_56 = cmd_alloc; CmdRef param_57 = cmd_ref; - CmdImage param_58 = _1850; + CmdImage param_58 = _1845; Cmd_Image_write(param_56, param_57, param_58); cmd_ref.offset += 12u; break; } case 5u: { - bool _1864 = tile_1.tile.offset == 0u; - bool _1870; - if (_1864) + bool _1859 = tile_1.tile.offset == 0u; + bool _1865; + if (_1859) { - _1870 = tile_1.backdrop == 0; + _1865 = tile_1.backdrop == 0; } else { - _1870 = _1864; + _1865 = _1859; } - if (_1870) + if (_1865) { clip_zero_depth = clip_depth + 1u; } @@ -1020,11 +1011,11 @@ void comp_main() Alloc param_59 = cmd_alloc; CmdRef param_60 = cmd_ref; uint param_61 = cmd_limit; - bool _1882 = alloc_cmd(param_59, param_60, param_61); + bool _1877 = alloc_cmd(param_59, param_60, param_61); cmd_alloc = param_59; cmd_ref = param_60; cmd_limit = param_61; - if (!_1882) + if (!_1877) { break; } @@ -1042,11 +1033,11 @@ void comp_main() Alloc param_64 = cmd_alloc; CmdRef param_65 = cmd_ref; uint param_66 = cmd_limit; - bool _1910 = alloc_cmd(param_64, param_65, param_66); + bool _1905 = alloc_cmd(param_64, param_65, param_66); cmd_alloc = param_64; cmd_ref = param_65; cmd_limit = param_66; - if (!_1910) + if (!_1905) { break; } @@ -1057,10 +1048,10 @@ void comp_main() write_fill(param_67, param_68, param_69, param_70); cmd_ref = param_68; uint blend_1 = _1222.Load(dd_1 * 4 + 0); - CmdEndClip _1933 = { blend_1 }; + CmdEndClip _1928 = { blend_1 }; Alloc param_71 = cmd_alloc; CmdRef param_72 = cmd_ref; - CmdEndClip param_73 = _1933; + CmdEndClip param_73 = _1928; Cmd_EndClip_write(param_71, param_72, param_73); cmd_ref.offset += 8u; break; @@ -1095,17 +1086,17 @@ void comp_main() break; } } - bool _1980 = (bin_tile_x + tile_x) < _854.Load(8); - bool _1989; - if (_1980) + bool _1975 = (bin_tile_x + tile_x) < _854.Load(8); + bool _1984; + if (_1975) { - _1989 = (bin_tile_y + tile_y) < _854.Load(12); + _1984 = (bin_tile_y + tile_y) < _854.Load(12); } else { - _1989 = _1980; + _1984 = _1975; } - if (_1989) + if (_1984) { Alloc param_74 = cmd_alloc; CmdRef param_75 = cmd_ref; diff --git a/piet-gpu/shader/gen/coarse.msl b/piet-gpu/shader/gen/coarse.msl index 29174f3..4226352 100644 --- a/piet-gpu/shader/gen/coarse.msl +++ b/piet-gpu/shader/gen/coarse.msl @@ -874,22 +874,13 @@ kernel void main0(constant uint* spvBufferSizeConstants [[buffer(25)]], device M { _1551 = _1542; } - bool _1558; - if (!_1551) - { - _1558 = is_clip && is_blend; - } - else - { - _1558 = _1551; - } - include_tile = _1558; + include_tile = _1551 || is_blend; } if (include_tile) { uint el_slice = el_ix / 32u; uint el_mask = 1u << (el_ix & 31u); - uint _1578 = atomic_fetch_or_explicit((threadgroup atomic_uint*)&sh_bitmaps[el_slice][(y * 16u) + x], el_mask, memory_order_relaxed); + uint _1573 = 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); @@ -934,11 +925,11 @@ kernel void main0(constant uint* spvBufferSizeConstants [[buffer(25)]], device M Alloc param_29 = cmd_alloc; CmdRef param_30 = cmd_ref; uint param_31 = cmd_limit; - bool _1702 = alloc_cmd(param_29, param_30, param_31, v_242, v_242BufferSize); + bool _1697 = alloc_cmd(param_29, param_30, param_31, v_242, v_242BufferSize); cmd_alloc = param_29; cmd_ref = param_30; cmd_limit = param_31; - if (!_1702) + if (!_1697) { break; } @@ -961,11 +952,11 @@ kernel void main0(constant uint* spvBufferSizeConstants [[buffer(25)]], device M Alloc param_39 = cmd_alloc; CmdRef param_40 = cmd_ref; uint param_41 = cmd_limit; - bool _1743 = alloc_cmd(param_39, param_40, param_41, v_242, v_242BufferSize); + bool _1738 = alloc_cmd(param_39, param_40, param_41, v_242, v_242BufferSize); cmd_alloc = param_39; cmd_ref = param_40; cmd_limit = param_41; - if (!_1743) + if (!_1738) { break; } @@ -993,11 +984,11 @@ kernel void main0(constant uint* spvBufferSizeConstants [[buffer(25)]], device M Alloc param_49 = cmd_alloc; CmdRef param_50 = cmd_ref; uint param_51 = cmd_limit; - bool _1811 = alloc_cmd(param_49, param_50, param_51, v_242, v_242BufferSize); + bool _1806 = alloc_cmd(param_49, param_50, param_51, v_242, v_242BufferSize); cmd_alloc = param_49; cmd_ref = param_50; cmd_limit = param_51; - if (!_1811) + if (!_1806) { break; } @@ -1019,17 +1010,17 @@ kernel void main0(constant uint* spvBufferSizeConstants [[buffer(25)]], device M } case 5u: { - bool _1864 = tile_1.tile.offset == 0u; - bool _1870; - if (_1864) + bool _1859 = tile_1.tile.offset == 0u; + bool _1865; + if (_1859) { - _1870 = tile_1.backdrop == 0; + _1865 = tile_1.backdrop == 0; } else { - _1870 = _1864; + _1865 = _1859; } - if (_1870) + if (_1865) { clip_zero_depth = clip_depth + 1u; } @@ -1038,11 +1029,11 @@ kernel void main0(constant uint* spvBufferSizeConstants [[buffer(25)]], device M Alloc param_59 = cmd_alloc; CmdRef param_60 = cmd_ref; uint param_61 = cmd_limit; - bool _1882 = alloc_cmd(param_59, param_60, param_61, v_242, v_242BufferSize); + bool _1877 = alloc_cmd(param_59, param_60, param_61, v_242, v_242BufferSize); cmd_alloc = param_59; cmd_ref = param_60; cmd_limit = param_61; - if (!_1882) + if (!_1877) { break; } @@ -1060,11 +1051,11 @@ kernel void main0(constant uint* spvBufferSizeConstants [[buffer(25)]], device M Alloc param_64 = cmd_alloc; CmdRef param_65 = cmd_ref; uint param_66 = cmd_limit; - bool _1910 = alloc_cmd(param_64, param_65, param_66, v_242, v_242BufferSize); + bool _1905 = alloc_cmd(param_64, param_65, param_66, v_242, v_242BufferSize); cmd_alloc = param_64; cmd_ref = param_65; cmd_limit = param_66; - if (!_1910) + if (!_1905) { break; } @@ -1112,17 +1103,17 @@ kernel void main0(constant uint* spvBufferSizeConstants [[buffer(25)]], device M break; } } - bool _1980 = (bin_tile_x + tile_x) < _854.conf.width_in_tiles; - bool _1989; - if (_1980) + bool _1975 = (bin_tile_x + tile_x) < _854.conf.width_in_tiles; + bool _1984; + if (_1975) { - _1989 = (bin_tile_y + tile_y) < _854.conf.height_in_tiles; + _1984 = (bin_tile_y + tile_y) < _854.conf.height_in_tiles; } else { - _1989 = _1980; + _1984 = _1975; } - if (_1989) + if (_1984) { Alloc param_74 = cmd_alloc; CmdRef param_75 = cmd_ref; diff --git a/piet-gpu/shader/gen/coarse.spv b/piet-gpu/shader/gen/coarse.spv index d246506f188134063483dd02d569b3f612a658a5..b85fd8cd546262ba315eca88fc6be6357bb82392 100644 GIT binary patch delta 10019 zcmZ9R3%pfTwZ<199Ms4N!ioyi!195bnVo#1D4HdjNvY}K00%q5IlPV#3NG+1<@K?} zd|Xnq#J3n}nUbi0_ON?@)zxm9)xBBverrar|9_wT!E)^%^UUuXW6Uwe8f(t6_Fi0g z{h({E8#HXtz}?!4VnDG&F|_#F{lg9#f9xBE6h+sZj)il&XLJ>O1{d2EZ|axzbS|6? zA6&cx-Zbdw{zPoviEMgzUr*;e*9<8pH?A7Lb9sDoWmivEZ(q|^Z&Q)YT+q3=uk*~N z_i6CH?m1hdb;9R$F6iuOMW0dQ3v0aZReVUX2(JI0xxJ0H@x#h<;R|PU^>%g4oi%G= zS6_d|gNxB zU*pf!_>bW;%AT1Wl&*QJU#=OvQsY}{d?=fa_4FRKpTTr+q}jgN(Q z_s;BcTf0}^4reRvmy+)R&n{B`p4$5N*7*C{it<%gFr{X2VvV0vFfUP1EPr$_R(A|4il^4y z-+t1N(ypfFAk3WYXjN`Ut8%+pl~1m5`&w0RPYZt@|I;cQ`4NRpEe6|TU{ebp1-Ge% zPlDUj!krg29yt8C@<$+BQ8cc!4TWC|w+)3m+lIomw+n@9Zx;$Tb-Pfw@$5n^UN(Vk zCPwn3x%7oT`1gnY>nH6 zl$UN_cfD;$f!u($p>PA*hQiN*(}uQk?d?L*Yi}0{ceV?KYi}3gU)c=k0^3m4z%~@F z-Zm8OJh#GiXcvlJd%IA$6|xJ3TcPDme%iJJY#WM!4sAo>I(~On8e9~&(c|R9;X5&uAIf)wr_l{>w4Xh?d_5A`%ScCp{rn9t z%11Di{k0g_#Cq3c6n`#d&Bb_^zmY%6Y<_yTTPchEo59*Y*|L8NTp7)=EH<`;Z$
sTju}ZB6oJhWb~iOe!?8v3k+yapRGmR|&a0Tv-nhyexIQH-5Ni z_tKuNpTye(J|C<8#CRmY&Fe*M)WvQuusXfj=jG{upa}fJBXpoz+$gF7oS8w$tS;r3M|T9cTBRf5%2K7Hz)5yjm6I3>%~ou--ZS}8LmIt{N10}E z41dfud;M6jG7ns0zaMPu*RWZSX3cTn0nEoT)GehzOvMNIqb_ISgJAtB_nQ2wpseM$ zET%vdPUO}Qm-@F}7c$-miB3QW&{B9vnVjlvV*!itPcoJBd=O{KGZrP+moD5gS zSD$T_R+t4J0pCwD;U9$?I9GKFSR3UYs6P?xRPbgJ_du!7+S9=5Onxe}Ji3or|1|{4 z{yQE1F-C4_2iTAMPu&@eRdiW`O1GB)tQ=vJ3JI zz(p$Ka2B|FTh4?ZOmk5ln~#H)hj)XkYtDlIhKeYU&1|sp79Um2fvh0AY`z}&?8f|a zN0iHurr>jts_A9KavoT@XEk>7;ijeT1g`%OE^Pr=x!1pt*-IApF{`r{ld(G+Ov7&4 zpgXl`G!3DOJzbS+_7}bsOswT(XL>$4QxEW^oRo_P{cI_Q70Tt4;8!w_#K%?4$H1>f1N7nlvEQx|KKpkHXoQw_N~rs9kIxrmXy z@Bp~gJjU2$m=HC&5X5v)^$QsAMX>#-Hr#r@-?5eO(k^2fM9x$in7N$pricT)D{^ zSe}621Z%I(!1Ba<+DE8?R2ukcaQ5}Lzy@_+|2MNAx3%*BVkmPPi>sATo&etluViGW zd=<&VmAjsF2FGcmWYJXiZHIOTl?EYH>cFF57>9$2o{tT?luTClV&*3Sz`mB8stZr$+YX`6X7Z8O8_c!S; z!3NJaZ@D&c@GEdN_^;v0s*k|mz!O}3*8Uc(u6e83f5q=03KRVIV1xUZv!rru;@}V9 zYVbe8m9;gGZ2$MUKfx1Reb)XNtS%b89n+Gh8gg~GD|P`_+tjY`v?+9l zZ8h8tg-V@_WtPXuIIvE#Nyo#Ly;L-&Upk@mRo^=pG2aucY$Hv_9QFb$TgJI}s~wGO zZ@BU&u#ryyE7K-^_JQjsYxV^z%bNYb$jauT6L929^>-oEnM$w4*N)2KbbqkdC)fVn zU~OE#1bzTm8|B)#_YVZ;j{DkDuFk#R-ge}+Tkbo|c7}m{f2p&4WjP4^2n%$eo%-~G zgB!yyA92PZNDgMC%!h)N?ZYmP z?&%2#b5Ey$<+-OH0_PN;2$m=Chr#J*CxPYZXD5SgXZmxU`oxKD1>>4gP4!T=RFxU#APWhu=Za7r^9EKg~A!6{7-Se`iZ+-cb* zb5YD=WS7hbyGwNFLiGu_0IWQlM4o^P!P;k&$P=&+teK6hGn$=P6ox=R)_`*D{j zU&K)6rVv~DY!c;3>|AgvaSm9X#Fl|m3I9_>-h8l4CXk8PG0f~d6%6rm<|T|g*^=eOj!dU>w?LU8pVcafbYPdSw_yjVkqPA*`UM|TNWnZH{vAdySK%GCQ1 zSP4!AE(6Qe`4IRNI2HIL_)~5B_XoKgqN6+n{sX+31v=1BeNMv_&2w`lk}DWFH&=m` zRe!f$4X>UXxpvi?=S`G%Fne=Xq0gmW3!aMH4RJl#&z7RCK+by}1cMQLnt2@ab<8@D z$H5KY>Os5_uB>{Fl$n73T(6(3Uk!GB_-Ek#FI@Q_0v``IA<#&LS>D3z$8D?bW`;5y zi@l*b@Ofs?=$sSevtWbf@ZSnoR!#cW!x2=oHv4KBEvJ0nR#~0pd$;&Lja{xDUEap3 z+XbWgmVA{4Q+nltm|54Kh{;`K^m@MZ8^>KEbKD9;a& z2O9GmJC_e4d4RFDmCKjFS}Qk~hnSO#x-T=78C#qS{R((VC4-0ID=Ha00@p@)GFb1< zo{EuQO8f~JXJ5q#<<_Jg1?xb$7oP;2^W9|w*aR(%_9o{AJkFd8(B?6QGB4n$W*r;LF2z^A;&^iWaoU`|A)z{LqO z(VQwX5H)8>5mTeNY_`ef$p)=#`v32{A0E&BvETl_wbov1?REB9=ibALtJ^QUrhQ2J zfE_xDqOI7X*sS=zO>uHG_*I8p}p)Qy$w~+11n4+t;+!+bPI8=S)7UZ}Rk}_bzx}_l!-^ zrom@To-?_p75%gtKfA`~zk&}g&V}p0XJ&7sW7Lpx5q$2{uHLRmGp9|P+tt^f@u1=& z#e<4#!85vhyUyzF>^rTM*D7SyJa4S=oBDaVX(M;k z3|80poi%<}jo%BOGkwbBl>FhEeqD_}+T`WSm3*>h@N|uDtnnYz_>XG**&6?8jXwvU zTK05KqIAtm{qLH=e`~zpFjh-Cq{d%U~k&^EQ&nZ&BZ*Be98h>v`QNChF$JGob*7%7vKBdMxYkXFX&#Cdgj`EcyEo%nn z*Z2iBeo>7thEJK%)yqxNtk7jO{n8p=KB6Z0QjI^gcH>T$w6|?sd;I7vw>_Xw% zOPajow79@F6ayF7hQgg~L*d%nhQhVC4TYPqT`1gy?Ly(sb|L37$JI8ZK<=rx4TT%f zHWY3^+fev9aN5u&uDx9-dhP8(;m&p;{*}#uF0c)$Faa0XhQig`hQig`hQgiuDqMR{ zcJ$iYg~F|nUC7qi3}}gLLosl+4TbB_HWaQy+few0@Z}Y*y}gYvaTD?<9!LcE#2Z18>f7fx-4CiH$hgy76S+uyPwDGE>_Y?8nrU zzmcKL)WoL3R1D+KV%A(tD)P7RN0}X1@3u^3u^$fB{)v|TTj9z^4CXqmZwa?Y@B+?M z>cl>)EKYU+>p1*v;Ky-#VQc+JxU!Qde{4pjocaOd6kg6z?Jn7!ELGArO|Wdj-?L{ z{UqLQ@Yz`PCsq`@!_Dj0*rPRcQRb4BXNSh8*H%8P)SeU7_c%o;L+I^>_ot=rm#2)*3G{xj z$>;EY0IbYAFg8bmP3Bc>ZvsJ&0=F?A%0O2Z3tBrr8imT-CC7kusN7xp4MtfU9|u-8 zjip@l@a@QZrEYJEAz@rD#p+C~PEl)(C&^s$eCpC>;?inY;m1&#O zD~tUH!D~n+{8YGsv#Ae(wNc*e)Wk0)K{k-Nx0nW5{9&*<1D(Pwk8U#fRs!WfPJt`S zHSPra@fxd}%2>%G!b>WSpK?0HG8$wWvs|UOS2uVaLDZcFmXF~&Zh`J}@OPM9qb?55 z09P-|4ETO@8|AU-0V{`>xGRcY{!|yug#UqxD9@rW7&e?ZRa-DEEqaqc3We^g7=L% zw_nSCCdO56dKC$Tl)RAMz) zuFiA67VJ4Uu{)XN>MZ@;U`wv$Owwql=zgtZ#gM(a`d7B+8GzMPSR@?CJ6 zhPtu-R%R}~#~+pH8s7&iGck{_JiGlNIOY8TSf1Vf7@YF{2rO6Ut@0CaZk1=j@=~Su z%FiGMGT={{<+)dW4)#_tm`+`+O@jWSB~OXv>X?dO@~3*M{0g311zowx@Yg6*=Ft5H z{2U{PPOeQH{}ya%b8vqLR^}I}l>7Jalv{n){sF9RLBYEccmGF-LW4gG{}b5Y`DQNH zCJz1#t_FV|u530>ss#TFJi*mx?O&BMnwOgMSNsj4Fv0&0Hn@*HODfkU4qgCPgTDw@ zR(<^b1D@dOv-Y20bvb9{zaR<|{3Wo#H|PF0xLg|@cyIp)oWu8Tuw0!_tp9;?`2Gu) ztLsJAM$6H!K5aj8Ztd-xZE67I9*lJ8LJe{b1~zWEbZ9vU$v{S~_F%9wUu?#*18n%5 z7`b-UoAatxN(R{+opsCw+X9}NsY}ha^p%Rr6zx@Dgw4Hp&&X4sSA$a@Ux3L|pVxv@ zA76&a6Xf;av_4;q$m)t+g-ux4C&iQd33vjm93}en8F@lWM#3C``)8A0pXqqRXm1~ z{GDKB+QiXba2;jM-e6@}^DeM5*Bpf(U;OU_*Oxkz^VO)bv35kWc{jvL7{mo*z#8db zKKwmkZIo+cZ`co^oL{AjD*n7|(oo%W7qZ5zv$3@h{7_At%(_!Zy~ z@SzN4-UebL-%N!M2kXg4hh8QyAAwAHemB1#tjtJe6hF?b_m{;7z{Dw=KT2F+5O2t% znEjZDMn^J^WtfP#nL7HT!M9L?l<63_GS97*_lLYP6VRvY_2YW4=JDXA?cDhq9|uvO z(#v@QIG1xGSf0Qqfz!%P1k2OPP6pe_^yfPD33v)vx$!(k^5z#L1AY+V0+j}qC*Y}I z9jP<0Jgw|QV1rt@4>S9*a>^$$l(}8TH|M8Y$DkKMc3Rtev#Ja$#L?>A8k8;0$ zPXn)Gc8$86&gso(+G$9pGjf5t!OHvv9?R3=)pIP@u6ohDcqzph;Oa%60aunUGU}AE zzTEUf>>7@19Q3vvtJBG1W_g@k3RdP%)C)-D<6vdk^5x)c`7*FP z$t~+R@a3z&6vZ-zL6$S?C>_56-hgu*Xs15Me?{{~xdO=wMsAcV!OE(?QLlnmZxp$9 z)vL9Vb1v^>_G+#~pRHaCJ_Wfa;FDlKFBKgH@{^l_K?pv9VkGl*%sPP_$)8OlxT7G^&lOm#Oil<7%qab52d%P7$~ zC(B!1gH3Ml+d7J43We1)5xyc#)2-Ic`pNfh)zMkLXN&LEShH+c`B_%oMyOo1+u_Rc zzd`DhCAH7B?8EQS9-BLOux`edJ`dMMd2CjLyRivh13$fg&BW%SI}vE4BL9217OXRM z*7a^?Kh{;tM|dRQJ!BH_cz)a+msI= zxu3DRwW$ZeS}Qk~hnbU$x`!Cbj4jSa9|6y=WUvl?ekFr1!nIMJ3?8*5OvK2q9sa(H zv&S$(xhbjjU>zuT^GU!tUr!zfo1lf!-sIfC6U^CwHX9hq+`xv8N&P{6a6E}Hz2i$@ kxe579eHxtR`4m{5XX=;1E3t8nIulHQC)Uneas5932j9!vEC2ui