From 2613a7e500a8f6c0a2ebf0a88302bb4558ce18c3 Mon Sep 17 00:00:00 2001 From: Raph Levien Date: Wed, 19 Jan 2022 12:18:32 -0800 Subject: [PATCH] Add generated kernel4_gray shaders --- piet-gpu/shader/gen/kernel4_gray.dxil | Bin 0 -> 9936 bytes piet-gpu/shader/gen/kernel4_gray.hlsl | 688 ++++++++++++++++++++++++ piet-gpu/shader/gen/kernel4_gray.msl | 727 ++++++++++++++++++++++++++ piet-gpu/shader/gen/kernel4_gray.spv | Bin 0 -> 38880 bytes 4 files changed, 1415 insertions(+) create mode 100644 piet-gpu/shader/gen/kernel4_gray.dxil create mode 100644 piet-gpu/shader/gen/kernel4_gray.hlsl create mode 100644 piet-gpu/shader/gen/kernel4_gray.msl create mode 100644 piet-gpu/shader/gen/kernel4_gray.spv diff --git a/piet-gpu/shader/gen/kernel4_gray.dxil b/piet-gpu/shader/gen/kernel4_gray.dxil new file mode 100644 index 0000000000000000000000000000000000000000..7bd557b984270e95bd0f3e3266d850a75a725d7b GIT binary patch literal 9936 zcmeHseN+=y*Y6~mOad7~!j}_1BoIMV)CnLWAWaA&Dk=&pimeGzL>nk7Qa>oo2PRR1 z1dHF+hN>;LwojBvYr&}iXBv3uVj{?6a2Dq; zAi)hmycYynP&{Ag8+wHz#X1c-`u|r<-_x@>u~Nr;VTg1KZ;AlB`1JFDX{7rU_@T&fcmkctn}{}*D~Wm3+yks! zJ0~S9gGnq3OexnF7{mCIeEk`Omy$bJz&3IP_9F2zm0>{h^*&UXk1Yevfc}XxxS`t^ zB5FsFb_5P-e9EU-xd#p@Mq~lN!F`9SrMdh31NiIrprL}(jo|XI zvE3LJQfM%+8tsJ~la(n7Kmv!D787nGIYL?yQ%>TU`((IQMzpnA34Bd; z%3ZVf`poL6aPY(yP077E*Y@>Nm!zP9nP$s1xu+~l71>B2q^LwMBd}w=RNP5IDeM*7 zP_fWt4KH7RW|0;N)D6mSYRlTQD+D=b-wek8_=DkjBT>PX> zr}nA!MWpfhV`=l+$gydcWufgslDvXj`QdbKyOdu;sK|cvy^+sZPy}-dFK`ri{RLexR8&Dd-DdHa(6COf%Ik-2XNg$z+9*%5EMK zzY%eqpqQQ;XY5Xp&w3y)l2&NUr-tUuBIK-XS+ve+FCAaBdmHWOmM)68ot1Tbck9&h zse#9){(7uSwNtog8%w_Ea}FBXC_Bo>`!E?lor#KNifT2P7Nw=dB{E|PqvulVQSkJp zwy`LpoFJ$!PFiRqIb2YG-Qp&E|;l!mxQ3RE1~H*iyHN{2Z8tOJr`eEaqr8*^Jm- zvx@JO7Hfj`5GN=c>{&}v;gT82KQw(;m2CKKQF6=MaBD@4-GhVOqaV8NU+Z>#=yD7V zCo(1T@(O44RpDomcj>f<;5X)4cni zMSppGEwAZP+U!eR0o|r5;}@2o`w07i`+#PSNt0>TWLY%XPK_+sV+xEq{`B!hk*mSJ z+xQ*1+HcBNCC)zl^H6#M|Cn*0%eANT^MltH2kgsSzhlO>O4<4iTNgjKeuJWzdK%}Q zY+C#Iy|S0@m41G&4Z)|j4s9@hjGa3e+m39(hmhd! zIPAwtb}F7bvD+GRvpu?kqzM)m+1TulZ)9AS?Xu~RwrnEP?juDFx5c|ARw~ndXLg==F!gWNI{QU^~dpG+h3;PjC zD|fU<54K15d9U|m(Y;`3;sB!FS)m=*3+7JK3L%D@;hXn6e1>UY1yZMk&vW&CrTSU`x z@4bWq+t*nwGxJcA8|TDXGV~?{1m9f$kE>N>BJaETgGDVF9!J%2@hm7Rbj#wByf!xJ zxmp#a{;w7DOQGC`Eq4=*Q2QUXK%idZQnpIw=4`#Yas&s9IiCU>CcIyTpfr9EFccgG z;^G2taR%eijqbdoAFSFR{`FGE>(kf$0{ivtnDPF-9rWve?%$Wey1#!bz_WkV{=E|H z>wnw7?f>1s|4;gNh?Sp4;bO}}IadB5!xSVE=F7iQKL_^sj(?ec&ISg7ejfJMejdHA zJ-OG!+;@r-sD=%EJ`xC8xPynJY($cLEgaTmIe7p7wSR+7c3w>2fpXA%t)G+p0HN5E zP6dtlb0dc(!6gzTin>wF!M;ET-bhbI?Hs!^Zy8$1aUjRXJ9Q#cb3x2M7xWP+L}!&O z3ws)T8+>6d36-pugwN}rB_=frTW%G;ENDg8yw9-Q?zC#!kF^LX<{%+evQ5>Yjj0zC znY4Fh9F4*y;>e86R_y~C%_5*hR4AU4)z;tt;@&Oy{o&qw!=h`>?m=SlnRl(72UYkL z{H=4!q5wL*is_W9Yp^@o0sNLIC>^gbW1Gz~DJ{K-#@j-rOk4Ym#tA9;_&qJ z57G~(y^}Kg(iCyg9{l60Zx>e2`0hs4LQz_rZCfWxeja=Pg%H~vBWs78-HzVDWA`1s zuGi27yp-$-UGH?gpV`&`TuXPl<^Uy z6a2UiMK^O;;74u>;AUDM5>_sXJGiW?ZK13!HzBOR+V-dDdidh@u;)Z$SlGOQure+d zRx-wD`}0rng4-eqYl0W>FjLf&#-dys$9DQf_ z=QP5=b;;j)?<+za+ce{kX70rolE3(n%y<4Wz3Jb(;cI`!0+mBo-hO9xKBDI9Ulf2{ zAcz19K-5y@(QaPpi*E4asoaS0nV4&e$@&xH1vA6)jto^{ zwv@nBZY7y~>i=%iJ}y831(YkH^5Zj3Se5$q#8uGhg||bcXJsFqwTVG+AQ`%QX%4$lZXwo^D!EjFujvJFfltR0^QcqVdX*1 zoa9(qRPq!3?zwjBBa(YH^Qi0TWh?1DSXWi00Wf znaQ<$+Rjp-b3GaS2rAk-z88jfFWyNWp8*vG4YyB-$TW%Pdb)ENX=dFzururz#Je3Y z8m}AQj(38B2`k2h#6ALwq6TeqlWV247c4-L!W=vX73~_|(YPY-B+ukk#P-NyBQh=G zr$)ezGDdU_U`LEU-NA52 z=6SS|JZ8cymDDICd_+~^W)!LBtPQ91g@m0)dqh|Hc)zen=yY&m5)5dGC^&OGE$%KF zu}wjLCy#NEaFd2T8zrGc6&K$Y)Xhz3qOJa97Ix&+te8Dc3h#oIht~!fx16%_E>}vi z9|?oiv&Z`KmcD~e-5D!+*iYI;pCLRa&n}l@-w_5%7{*~dZKSU#qH)ioSCdslTX#qi z<*7>gbc$~?19=N9xzdnR0hFS~0*jtP1r$(ZsY1?kPTRo-6}1_AUUu!3(c&w06xXma zFuOn0hUK~0G;@Dk_%7;XZIIp_gGo$D+#AxX4w+Z?iVIb~5LeXifU_IZny z5^1Q{yh6Wu&#FQnat5J!x8p+n*9%A>tP93pZ+j$=SNOaXbqjzHRYpgg*>?ibAWlO;P`ofF{#q?*^(zmy7 z|Bb%VG~s#8vdP2-JrXNHJqBVDkT33zl>L~I>nlAAp4HPIi&_FGFXZEF1hu-AL6?{d z%k6GN#KHHw%uVh|$)-U$H6b?ifai$O-h+yB5}BjIIOWIVl_&Xq*#v98u0bEvOqCmm zXt)DwR67?3r03wu?e9`F16tE8i}a9cwYncoL9|ioyNG7BHb3BIu;E=%vRm?^6c1Gc z>%8wSn|PKxY}Z?1Rf`_#jHE18)D{GMHPQ^C$-$KQR-z95ktfVRIH*Q4cdfv)m#d*6 zVlv|nEnRJP4ZA}giYJ0#^Ks;5wa^goAho;(*0G-5efgpY0mcAisIkniLMuA^Wr&%F3jHrQbk1~j{~of)AS$vY#i&4 zQkD6PZ;$;quH|aJ?xeBhCy+^08B!Txkd<*>@!9yBBwO7@j2r*(DTkn=shM5ACDsiS zts8MpCu|n}6TyP56DiO51dHRp`3JDL ztdXf7xAQA+|m<@(}V=Tzt8kH<6 z=6m(L()A$s8zGY>On(p$oJ&HwT&`&ij9)3X7N<5AdDi&ymftg129gUf}BS`7F~5lv<;OwW7nvf{Ct^H9^c75;Q)LI z6UO`a64i2=S%pqb0esnRVEQ)kvmkkEl(-3QVoDs{}%+<~TVZRz}s*xl55nAoiY+;WPpDiLiwg^bMH`pNi~ zZcZW^chN%9LbLQO^yd)2l`i^9HRTSWT!du3#yGZ7S(43s8%AT@AlkoBeC zUwvBb^ZDbiS5Y%SLiWZHNIV17BKW~%eRbxKAwClI&wLQlcil|c!nPvDIfUG9fkp06 zL;GQ$SJ}oXhb>D$>eezbrPE4`e`lGMeOFMQ2MzcobZ3e}bWTDdr`gvCAqnfyhV_pV ze8IM@zmwp8Nk)4$Mor<@k=J}hGVzm{R$slC@^FUr72t&*I5^U!5$5AKD7(@HJY#^5 zzS5&R0`fMvlzN>M+5FT!dPTtgW_Hc;l!R2kqcbU>9X+tGUs3@mS-1gE&vE^_z7J6E zL@%J&kE8CZpR6GM~-3rLQm(LaU0voJvKq6a^o^B6)E&AYJ)Z|^!I@W6Q@2qr1B=X1ZSla-16{~{ z*rkCXW5Bz%g$s zZOv2^`-bYwDv57+AhSBvRg^l~X{{KP+io~E`qJvo-s7_IVZ>*x{^6m|jvkj?k{RFl zRx5tvwJ7ZDu>=P;6+2;jzz6kU=3AjfQk;d49g#YEFiB1#)>t2rd7|%8G%IF_w$8YR zJ-ImRdamxJ&X_$+IY)-a8^n#m!i}JE6gheV#W^eS(N49oh0Cos=-u)>G*%D~3C9(J zPiZs_x0oqEGU2yth7}=J=8!}+a~A`z;OUb)o*9ejaBsgr9Z~s$euwe=wM56uAm~Z2 z)=BLM1yCjZm(OF)2kh9G7Ki-3BZe4{#q@^y$yMsFU@VqbF!FH@Oy^UtjL43Q>hnXX zZnpey#2jr;|f!gn@C$>F9lld{MCG&H9)Vyf?6`g>ml{;aJ?`Iuc`t9J-k*}BjF}U;( z&<`w?YK|ch@2;c@@%fy$Q|oP9`@iE~)u5Ct+s64nIC|J$fdcq%g4p5)M$T)>;QGRV zSdWX7GukPx2eI@#i6xoR?vJG~(?2{U%x$-7e`#$A8WF@sJ$Pc%iT=PKpLr%O%RfB5 zj1}}rkfvnuY<&5wFUf*Wa~WR>Pt7hJHf3!@8LXIEM)Lz{yXO1TtaIg){}Lw_-iR<>?PDn@>Unkq80_@^&-x;I$UN0QjOa``YsG3nL561C?OwY5evvPw#RL}KNf zw%9WfVM%)<`a+`V3uW~4=mxS#=u(KV*~_qTrR3E+s{<-lh&L?=q9kv3y#85W5|K}@ zJCo26u~NBejr|!Jeb<={BjG1yx5M*eWVgBbyv~)QLiJmq?Uh~cVZMLht6pnASYOyWZH!1=HQKr1?LoP9t+nUGr2{SW?ey5D?-%Ao+zYSwdB62< zcc=P>M`zyi$E5vl_aYU@^f4D_4Jeft22*`sYYji%@%_naQNvIEVN=7{9TFD2S{r`_ ztbo&NOW*uQxn>Xh**)v3%ud-`qaOuDWS%?NaD~FX=FL_OQ?95jdh^d$X@LfrDlRq@ zA9I0@$du$9Wc;4+E(3U6`DQ39iPXla>T&L%+`*5PFvo~z=|6UgUTO(ARZMetvsc}sa zb+c8kqJf{VC7Z9UdzgTx*G2bqt~+#s2g^?^4V0fWkp`iaBh}zj5NwHs;N_EmWZOOnu55+inITYe zJ_Pp_KyYLL0M`g0cwqvx`62{!4ny#$7rYLFM;1YF0|X-}5S*I`!Pg-;=_mvzg+j?2 zAlO_8!H2xyR0uBjf^R@@hZmdy!K#fA?15m+48i<#2zEfQ=5q)Z9fy)PL$ItGf?bQD z image_atlas : register(u3, space0); +RWTexture2D gradients : register(u4, space0); +RWTexture2D image : register(u2, space0); + +static uint3 gl_WorkGroupID; +static uint3 gl_LocalInvocationID; +struct SPIRV_Cross_Input +{ + uint3 gl_WorkGroupID : SV_GroupID; + uint3 gl_LocalInvocationID : SV_GroupThreadID; +}; + +uint spvPackUnorm4x8(float4 value) +{ + uint4 Packed = uint4(round(saturate(value) * 255.0)); + return Packed.x | (Packed.y << 8) | (Packed.z << 16) | (Packed.w << 24); +} + +float4 spvUnpackUnorm4x8(uint value) +{ + uint4 Packed = uint4(value & 0xff, (value >> 8) & 0xff, (value >> 16) & 0xff, value >> 24); + return float4(Packed) / 255.0; +} + +Alloc slice_mem(Alloc a, uint offset, uint size) +{ + Alloc _215 = { a.offset + offset }; + return _215; +} + +bool touch_mem(Alloc alloc, uint offset) +{ + return true; +} + +uint read_mem(Alloc alloc, uint offset) +{ + Alloc param = alloc; + uint param_1 = offset; + if (!touch_mem(param, param_1)) + { + return 0u; + } + uint v = _202.Load(offset * 4 + 8); + return v; +} + +CmdTag Cmd_tag(Alloc a, CmdRef ref) +{ + Alloc param = a; + uint param_1 = ref.offset >> uint(2); + uint tag_and_flags = read_mem(param, param_1); + CmdTag _432 = { tag_and_flags & 65535u, tag_and_flags >> uint(16) }; + return _432; +} + +CmdStroke CmdStroke_read(Alloc a, CmdStrokeRef ref) +{ + uint ix = ref.offset >> uint(2); + Alloc param = a; + uint param_1 = ix + 0u; + uint raw0 = read_mem(param, param_1); + Alloc param_2 = a; + uint param_3 = ix + 1u; + uint raw1 = read_mem(param_2, param_3); + CmdStroke s; + s.tile_ref = raw0; + s.half_width = asfloat(raw1); + return s; +} + +CmdStroke Cmd_Stroke_read(Alloc a, CmdRef ref) +{ + CmdStrokeRef _449 = { ref.offset + 4u }; + Alloc param = a; + CmdStrokeRef param_1 = _449; + return CmdStroke_read(param, param_1); +} + +Alloc new_alloc(uint offset, uint size, bool mem_ok) +{ + Alloc a; + a.offset = offset; + return a; +} + +TileSeg TileSeg_read(Alloc a, TileSegRef ref) +{ + uint ix = ref.offset >> uint(2); + Alloc param = a; + uint param_1 = ix + 0u; + uint raw0 = read_mem(param, param_1); + Alloc param_2 = a; + uint param_3 = ix + 1u; + uint raw1 = read_mem(param_2, param_3); + Alloc param_4 = a; + uint param_5 = ix + 2u; + uint raw2 = read_mem(param_4, param_5); + Alloc param_6 = a; + uint param_7 = ix + 3u; + uint raw3 = read_mem(param_6, param_7); + Alloc param_8 = a; + uint param_9 = ix + 4u; + uint raw4 = read_mem(param_8, param_9); + Alloc param_10 = a; + uint param_11 = ix + 5u; + uint raw5 = read_mem(param_10, param_11); + TileSeg s; + s.origin = float2(asfloat(raw0), asfloat(raw1)); + s._vector = float2(asfloat(raw2), asfloat(raw3)); + s.y_edge = asfloat(raw4); + TileSegRef _572 = { raw5 }; + s.next = _572; + return s; +} + +uint2 chunk_offset(uint i) +{ + return uint2((i % 2u) * 8u, (i / 2u) * 4u); +} + +CmdFill CmdFill_read(Alloc a, CmdFillRef ref) +{ + uint ix = ref.offset >> uint(2); + Alloc param = a; + uint param_1 = ix + 0u; + uint raw0 = read_mem(param, param_1); + Alloc param_2 = a; + uint param_3 = ix + 1u; + uint raw1 = read_mem(param_2, param_3); + CmdFill s; + s.tile_ref = raw0; + s.backdrop = int(raw1); + return s; +} + +CmdFill Cmd_Fill_read(Alloc a, CmdRef ref) +{ + CmdFillRef _439 = { ref.offset + 4u }; + Alloc param = a; + CmdFillRef param_1 = _439; + return CmdFill_read(param, param_1); +} + +CmdAlpha CmdAlpha_read(Alloc a, CmdAlphaRef ref) +{ + uint ix = ref.offset >> uint(2); + Alloc param = a; + uint param_1 = ix + 0u; + uint raw0 = read_mem(param, param_1); + CmdAlpha s; + s.alpha = asfloat(raw0); + return s; +} + +CmdAlpha Cmd_Alpha_read(Alloc a, CmdRef ref) +{ + CmdAlphaRef _459 = { ref.offset + 4u }; + Alloc param = a; + CmdAlphaRef param_1 = _459; + return CmdAlpha_read(param, param_1); +} + +CmdColor CmdColor_read(Alloc a, CmdColorRef ref) +{ + uint ix = ref.offset >> uint(2); + Alloc param = a; + uint param_1 = ix + 0u; + uint raw0 = read_mem(param, param_1); + CmdColor s; + s.rgba_color = raw0; + return s; +} + +CmdColor Cmd_Color_read(Alloc a, CmdRef ref) +{ + CmdColorRef _469 = { ref.offset + 4u }; + Alloc param = a; + CmdColorRef param_1 = _469; + return CmdColor_read(param, param_1); +} + +float3 fromsRGB(float3 srgb) +{ + bool3 cutoff = bool3(srgb.x >= 0.040449999272823333740234375f.xxx.x, srgb.y >= 0.040449999272823333740234375f.xxx.y, srgb.z >= 0.040449999272823333740234375f.xxx.z); + float3 below = srgb / 12.9200000762939453125f.xxx; + float3 above = pow((srgb + 0.054999999701976776123046875f.xxx) / 1.05499994754791259765625f.xxx, 2.400000095367431640625f.xxx); + return float3(cutoff.x ? above.x : below.x, cutoff.y ? above.y : below.y, cutoff.z ? above.z : below.z); +} + +float4 unpacksRGB(uint srgba) +{ + float4 color = spvUnpackUnorm4x8(srgba).wzyx; + float3 param = color.xyz; + return float4(fromsRGB(param), color.w); +} + +CmdLinGrad CmdLinGrad_read(Alloc a, CmdLinGradRef ref) +{ + uint ix = ref.offset >> uint(2); + Alloc param = a; + uint param_1 = ix + 0u; + uint raw0 = read_mem(param, param_1); + Alloc param_2 = a; + uint param_3 = ix + 1u; + uint raw1 = read_mem(param_2, param_3); + Alloc param_4 = a; + uint param_5 = ix + 2u; + uint raw2 = read_mem(param_4, param_5); + Alloc param_6 = a; + uint param_7 = ix + 3u; + uint raw3 = read_mem(param_6, param_7); + CmdLinGrad s; + s.index = raw0; + s.line_x = asfloat(raw1); + s.line_y = asfloat(raw2); + s.line_c = asfloat(raw3); + return s; +} + +CmdLinGrad Cmd_LinGrad_read(Alloc a, CmdRef ref) +{ + CmdLinGradRef _479 = { ref.offset + 4u }; + Alloc param = a; + CmdLinGradRef param_1 = _479; + return CmdLinGrad_read(param, param_1); +} + +CmdImage CmdImage_read(Alloc a, CmdImageRef ref) +{ + uint ix = ref.offset >> uint(2); + Alloc param = a; + uint param_1 = ix + 0u; + uint raw0 = read_mem(param, param_1); + Alloc param_2 = a; + uint param_3 = ix + 1u; + uint raw1 = read_mem(param_2, param_3); + CmdImage s; + s.index = raw0; + s.offset = int2(int(raw1 << uint(16)) >> 16, int(raw1) >> 16); + return s; +} + +CmdImage Cmd_Image_read(Alloc a, CmdRef ref) +{ + CmdImageRef _489 = { ref.offset + 4u }; + Alloc param = a; + CmdImageRef param_1 = _489; + return CmdImage_read(param, param_1); +} + +void fillImage(out float4 spvReturnValue[8], uint2 xy, CmdImage cmd_img) +{ + float4 rgba[8]; + for (uint i = 0u; i < 8u; i++) + { + uint param = i; + int2 uv = int2(xy + chunk_offset(param)) + cmd_img.offset; + float4 fg_rgba = image_atlas[uv]; + float3 param_1 = fg_rgba.xyz; + float3 _695 = fromsRGB(param_1); + fg_rgba.x = _695.x; + fg_rgba.y = _695.y; + fg_rgba.z = _695.z; + rgba[i] = fg_rgba; + } + spvReturnValue = rgba; +} + +float3 tosRGB(float3 rgb) +{ + bool3 cutoff = bool3(rgb.x >= 0.003130800090730190277099609375f.xxx.x, rgb.y >= 0.003130800090730190277099609375f.xxx.y, rgb.z >= 0.003130800090730190277099609375f.xxx.z); + float3 below = 12.9200000762939453125f.xxx * rgb; + float3 above = (1.05499994754791259765625f.xxx * pow(rgb, 0.416660010814666748046875f.xxx)) - 0.054999999701976776123046875f.xxx; + return float3(cutoff.x ? above.x : below.x, cutoff.y ? above.y : below.y, cutoff.z ? above.z : below.z); +} + +uint packsRGB(inout float4 rgba) +{ + float3 param = rgba.xyz; + rgba = float4(tosRGB(param), rgba.w); + return spvPackUnorm4x8(rgba.wzyx); +} + +CmdJump CmdJump_read(Alloc a, CmdJumpRef ref) +{ + uint ix = ref.offset >> uint(2); + Alloc param = a; + uint param_1 = ix + 0u; + uint raw0 = read_mem(param, param_1); + CmdJump s; + s.new_ref = raw0; + return s; +} + +CmdJump Cmd_Jump_read(Alloc a, CmdRef ref) +{ + CmdJumpRef _499 = { ref.offset + 4u }; + Alloc param = a; + CmdJumpRef param_1 = _499; + return CmdJump_read(param, param_1); +} + +void comp_main() +{ + uint tile_ix = (gl_WorkGroupID.y * _723.Load(8)) + gl_WorkGroupID.x; + Alloc _738; + _738.offset = _723.Load(24); + Alloc param; + param.offset = _738.offset; + uint param_1 = tile_ix * 1024u; + uint param_2 = 1024u; + Alloc cmd_alloc = slice_mem(param, param_1, param_2); + CmdRef _747 = { cmd_alloc.offset }; + CmdRef cmd_ref = _747; + uint2 xy_uint = uint2(gl_LocalInvocationID.x + (16u * gl_WorkGroupID.x), gl_LocalInvocationID.y + (16u * gl_WorkGroupID.y)); + float2 xy = float2(xy_uint); + float4 rgba[8]; + for (uint i = 0u; i < 8u; i++) + { + rgba[i] = 0.0f.xxxx; + } + uint clip_depth = 0u; + bool mem_ok = _202.Load(4) == 0u; + float df[8]; + TileSegRef tile_seg_ref; + float area[8]; + uint blend_stack[128][8]; + float blend_alpha_stack[128][8]; + while (mem_ok) + { + Alloc param_3 = cmd_alloc; + CmdRef param_4 = cmd_ref; + uint tag = Cmd_tag(param_3, param_4).tag; + if (tag == 0u) + { + break; + } + switch (tag) + { + case 2u: + { + Alloc param_5 = cmd_alloc; + CmdRef param_6 = cmd_ref; + CmdStroke stroke = Cmd_Stroke_read(param_5, param_6); + for (uint k = 0u; k < 8u; k++) + { + df[k] = 1000000000.0f; + } + TileSegRef _842 = { stroke.tile_ref }; + tile_seg_ref = _842; + do + { + uint param_7 = tile_seg_ref.offset; + uint param_8 = 24u; + bool param_9 = mem_ok; + Alloc param_10 = new_alloc(param_7, param_8, param_9); + TileSegRef param_11 = tile_seg_ref; + TileSeg seg = TileSeg_read(param_10, param_11); + float2 line_vec = seg._vector; + for (uint k_1 = 0u; k_1 < 8u; k_1++) + { + float2 dpos = (xy + 0.5f.xx) - seg.origin; + uint param_12 = k_1; + dpos += float2(chunk_offset(param_12)); + float t = clamp(dot(line_vec, dpos) / dot(line_vec, line_vec), 0.0f, 1.0f); + df[k_1] = min(df[k_1], length((line_vec * t) - dpos)); + } + tile_seg_ref = seg.next; + } while (tile_seg_ref.offset != 0u); + for (uint k_2 = 0u; k_2 < 8u; k_2++) + { + area[k_2] = clamp((stroke.half_width + 0.5f) - df[k_2], 0.0f, 1.0f); + } + cmd_ref.offset += 12u; + break; + } + case 1u: + { + Alloc param_13 = cmd_alloc; + CmdRef param_14 = cmd_ref; + CmdFill fill = Cmd_Fill_read(param_13, param_14); + for (uint k_3 = 0u; k_3 < 8u; k_3++) + { + area[k_3] = float(fill.backdrop); + } + TileSegRef _964 = { fill.tile_ref }; + tile_seg_ref = _964; + do + { + uint param_15 = tile_seg_ref.offset; + uint param_16 = 24u; + bool param_17 = mem_ok; + Alloc param_18 = new_alloc(param_15, param_16, param_17); + TileSegRef param_19 = tile_seg_ref; + TileSeg seg_1 = TileSeg_read(param_18, param_19); + for (uint k_4 = 0u; k_4 < 8u; k_4++) + { + uint param_20 = k_4; + float2 my_xy = xy + float2(chunk_offset(param_20)); + float2 start = seg_1.origin - my_xy; + float2 end = start + seg_1._vector; + float2 window = clamp(float2(start.y, end.y), 0.0f.xx, 1.0f.xx); + if (window.x != window.y) + { + float2 t_1 = (window - start.y.xx) / seg_1._vector.y.xx; + float2 xs = float2(lerp(start.x, end.x, t_1.x), lerp(start.x, end.x, t_1.y)); + float xmin = min(min(xs.x, xs.y), 1.0f) - 9.9999999747524270787835121154785e-07f; + float xmax = max(xs.x, xs.y); + float b = min(xmax, 1.0f); + float c = max(b, 0.0f); + float d = max(xmin, 0.0f); + float a = ((b + (0.5f * ((d * d) - (c * c)))) - xmin) / (xmax - xmin); + area[k_4] += (a * (window.x - window.y)); + } + area[k_4] += (sign(seg_1._vector.x) * clamp((my_xy.y - seg_1.y_edge) + 1.0f, 0.0f, 1.0f)); + } + tile_seg_ref = seg_1.next; + } while (tile_seg_ref.offset != 0u); + for (uint k_5 = 0u; k_5 < 8u; k_5++) + { + area[k_5] = min(abs(area[k_5]), 1.0f); + } + cmd_ref.offset += 12u; + break; + } + case 3u: + { + for (uint k_6 = 0u; k_6 < 8u; k_6++) + { + area[k_6] = 1.0f; + } + cmd_ref.offset += 4u; + break; + } + case 4u: + { + Alloc param_21 = cmd_alloc; + CmdRef param_22 = cmd_ref; + CmdAlpha alpha = Cmd_Alpha_read(param_21, param_22); + for (uint k_7 = 0u; k_7 < 8u; k_7++) + { + area[k_7] = alpha.alpha; + } + cmd_ref.offset += 8u; + break; + } + case 5u: + { + Alloc param_23 = cmd_alloc; + CmdRef param_24 = cmd_ref; + CmdColor color = Cmd_Color_read(param_23, param_24); + uint param_25 = color.rgba_color; + float4 fg = unpacksRGB(param_25); + for (uint k_8 = 0u; k_8 < 8u; k_8++) + { + float4 fg_k = fg * area[k_8]; + rgba[k_8] = (rgba[k_8] * (1.0f - fg_k.w)) + fg_k; + } + cmd_ref.offset += 8u; + break; + } + case 6u: + { + Alloc param_26 = cmd_alloc; + CmdRef param_27 = cmd_ref; + CmdLinGrad lin = Cmd_LinGrad_read(param_26, param_27); + float d_1 = ((lin.line_x * xy.x) + (lin.line_y * xy.y)) + lin.line_c; + for (uint k_9 = 0u; k_9 < 8u; k_9++) + { + uint param_28 = k_9; + float2 chunk_xy = float2(chunk_offset(param_28)); + float my_d = (d_1 + (lin.line_x * chunk_xy.x)) + (lin.line_y * chunk_xy.y); + int x = int(round(clamp(my_d, 0.0f, 1.0f) * 511.0f)); + float4 fg_rgba = gradients[int2(x, int(lin.index))]; + float3 param_29 = fg_rgba.xyz; + float3 _1298 = fromsRGB(param_29); + fg_rgba.x = _1298.x; + fg_rgba.y = _1298.y; + fg_rgba.z = _1298.z; + rgba[k_9] = fg_rgba; + } + cmd_ref.offset += 20u; + break; + } + case 7u: + { + Alloc param_30 = cmd_alloc; + CmdRef param_31 = cmd_ref; + CmdImage fill_img = Cmd_Image_read(param_30, param_31); + uint2 param_32 = xy_uint; + CmdImage param_33 = fill_img; + float4 _1327[8]; + fillImage(_1327, param_32, param_33); + float4 img[8] = _1327; + for (uint k_10 = 0u; k_10 < 8u; k_10++) + { + float4 fg_k_1 = img[k_10] * area[k_10]; + rgba[k_10] = (rgba[k_10] * (1.0f - fg_k_1.w)) + fg_k_1; + } + cmd_ref.offset += 12u; + break; + } + case 8u: + { + for (uint k_11 = 0u; k_11 < 8u; k_11++) + { + uint d_2 = min(clip_depth, 127u); + float4 param_34 = float4(rgba[k_11]); + uint _1390 = packsRGB(param_34); + blend_stack[d_2][k_11] = _1390; + blend_alpha_stack[d_2][k_11] = clamp(abs(area[k_11]), 0.0f, 1.0f); + rgba[k_11] = 0.0f.xxxx; + } + clip_depth++; + cmd_ref.offset += 4u; + break; + } + case 9u: + { + clip_depth--; + for (uint k_12 = 0u; k_12 < 8u; k_12++) + { + uint d_3 = min(clip_depth, 127u); + uint param_35 = blend_stack[d_3][k_12]; + float4 bg = unpacksRGB(param_35); + float4 fg_1 = (rgba[k_12] * area[k_12]) * blend_alpha_stack[d_3][k_12]; + rgba[k_12] = (bg * (1.0f - fg_1.w)) + fg_1; + } + cmd_ref.offset += 4u; + break; + } + case 10u: + { + Alloc param_36 = cmd_alloc; + CmdRef param_37 = cmd_ref; + CmdRef _1469 = { Cmd_Jump_read(param_36, param_37).new_ref }; + cmd_ref = _1469; + cmd_alloc.offset = cmd_ref.offset; + break; + } + } + } + for (uint i_1 = 0u; i_1 < 8u; i_1++) + { + uint param_38 = i_1; + image[int2(xy_uint + chunk_offset(param_38))] = rgba[i_1].w.x; + } +} + +[numthreads(8, 4, 1)] +void main(SPIRV_Cross_Input stage_input) +{ + gl_WorkGroupID = stage_input.gl_WorkGroupID; + gl_LocalInvocationID = stage_input.gl_LocalInvocationID; + comp_main(); +} diff --git a/piet-gpu/shader/gen/kernel4_gray.msl b/piet-gpu/shader/gen/kernel4_gray.msl new file mode 100644 index 0000000..e672020 --- /dev/null +++ b/piet-gpu/shader/gen/kernel4_gray.msl @@ -0,0 +1,727 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" +#pragma clang diagnostic ignored "-Wmissing-braces" + +#include +#include + +using namespace metal; + +template +struct spvUnsafeArray +{ + T elements[Num ? Num : 1]; + + thread T& operator [] (size_t pos) thread + { + return elements[pos]; + } + constexpr const thread T& operator [] (size_t pos) const thread + { + return elements[pos]; + } + + device T& operator [] (size_t pos) device + { + return elements[pos]; + } + constexpr const device T& operator [] (size_t pos) const device + { + return elements[pos]; + } + + constexpr const constant T& operator [] (size_t pos) const constant + { + return elements[pos]; + } + + threadgroup T& operator [] (size_t pos) threadgroup + { + return elements[pos]; + } + constexpr const threadgroup T& operator [] (size_t pos) const threadgroup + { + return elements[pos]; + } +}; + +struct Alloc +{ + uint offset; +}; + +struct CmdStrokeRef +{ + uint offset; +}; + +struct CmdStroke +{ + uint tile_ref; + float half_width; +}; + +struct CmdFillRef +{ + uint offset; +}; + +struct CmdFill +{ + uint tile_ref; + int backdrop; +}; + +struct CmdColorRef +{ + uint offset; +}; + +struct CmdColor +{ + uint rgba_color; +}; + +struct CmdLinGradRef +{ + uint offset; +}; + +struct CmdLinGrad +{ + uint index; + float line_x; + float line_y; + float line_c; +}; + +struct CmdImageRef +{ + uint offset; +}; + +struct CmdImage +{ + uint index; + int2 offset; +}; + +struct CmdAlphaRef +{ + uint offset; +}; + +struct CmdAlpha +{ + float alpha; +}; + +struct CmdJumpRef +{ + uint offset; +}; + +struct CmdJump +{ + uint new_ref; +}; + +struct CmdRef +{ + uint offset; +}; + +struct CmdTag +{ + uint tag; + uint flags; +}; + +struct TileSegRef +{ + uint offset; +}; + +struct TileSeg +{ + float2 origin; + float2 vector; + float y_edge; + TileSegRef next; +}; + +struct Memory +{ + uint mem_offset; + uint mem_error; + uint memory[1]; +}; + +struct Alloc_1 +{ + uint offset; +}; + +struct Config +{ + uint n_elements; + uint n_pathseg; + uint width_in_tiles; + uint height_in_tiles; + Alloc_1 tile_alloc; + Alloc_1 bin_alloc; + Alloc_1 ptcl_alloc; + Alloc_1 pathseg_alloc; + Alloc_1 anno_alloc; + Alloc_1 trans_alloc; + Alloc_1 bbox_alloc; + Alloc_1 drawmonoid_alloc; + uint n_trans; + uint n_path; + uint trans_offset; + uint linewidth_offset; + uint pathtag_offset; + uint pathseg_offset; +}; + +struct ConfigBuf +{ + Config conf; +}; + +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(8u, 4u, 1u); + +static inline __attribute__((always_inline)) +Alloc slice_mem(thread const Alloc& a, thread const uint& offset, thread const uint& size) +{ + return Alloc{ a.offset + offset }; +} + +static inline __attribute__((always_inline)) +bool touch_mem(thread const Alloc& alloc, thread const uint& offset) +{ + return true; +} + +static inline __attribute__((always_inline)) +uint read_mem(thread const Alloc& alloc, thread const uint& offset, device Memory& v_202) +{ + Alloc param = alloc; + uint param_1 = offset; + if (!touch_mem(param, param_1)) + { + return 0u; + } + uint v = v_202.memory[offset]; + return v; +} + +static inline __attribute__((always_inline)) +CmdTag Cmd_tag(thread const Alloc& a, thread const CmdRef& ref, device Memory& v_202) +{ + Alloc param = a; + uint param_1 = ref.offset >> uint(2); + uint tag_and_flags = read_mem(param, param_1, v_202); + return CmdTag{ tag_and_flags & 65535u, tag_and_flags >> uint(16) }; +} + +static inline __attribute__((always_inline)) +CmdStroke CmdStroke_read(thread const Alloc& a, thread const CmdStrokeRef& ref, device Memory& v_202) +{ + uint ix = ref.offset >> uint(2); + Alloc param = a; + uint param_1 = ix + 0u; + uint raw0 = read_mem(param, param_1, v_202); + Alloc param_2 = a; + uint param_3 = ix + 1u; + uint raw1 = read_mem(param_2, param_3, v_202); + CmdStroke s; + s.tile_ref = raw0; + s.half_width = as_type(raw1); + return s; +} + +static inline __attribute__((always_inline)) +CmdStroke Cmd_Stroke_read(thread const Alloc& a, thread const CmdRef& ref, device Memory& v_202) +{ + Alloc param = a; + CmdStrokeRef param_1 = CmdStrokeRef{ ref.offset + 4u }; + return CmdStroke_read(param, param_1, v_202); +} + +static inline __attribute__((always_inline)) +Alloc new_alloc(thread const uint& offset, thread const uint& size, thread const bool& mem_ok) +{ + Alloc a; + a.offset = offset; + return a; +} + +static inline __attribute__((always_inline)) +TileSeg TileSeg_read(thread const Alloc& a, thread const TileSegRef& ref, device Memory& v_202) +{ + uint ix = ref.offset >> uint(2); + Alloc param = a; + uint param_1 = ix + 0u; + uint raw0 = read_mem(param, param_1, v_202); + Alloc param_2 = a; + uint param_3 = ix + 1u; + uint raw1 = read_mem(param_2, param_3, v_202); + Alloc param_4 = a; + uint param_5 = ix + 2u; + uint raw2 = read_mem(param_4, param_5, v_202); + Alloc param_6 = a; + uint param_7 = ix + 3u; + uint raw3 = read_mem(param_6, param_7, v_202); + Alloc param_8 = a; + uint param_9 = ix + 4u; + uint raw4 = read_mem(param_8, param_9, v_202); + Alloc param_10 = a; + uint param_11 = ix + 5u; + uint raw5 = read_mem(param_10, param_11, v_202); + TileSeg s; + s.origin = float2(as_type(raw0), as_type(raw1)); + s.vector = float2(as_type(raw2), as_type(raw3)); + s.y_edge = as_type(raw4); + s.next = TileSegRef{ raw5 }; + return s; +} + +static inline __attribute__((always_inline)) +uint2 chunk_offset(thread const uint& i) +{ + return uint2((i % 2u) * 8u, (i / 2u) * 4u); +} + +static inline __attribute__((always_inline)) +CmdFill CmdFill_read(thread const Alloc& a, thread const CmdFillRef& ref, device Memory& v_202) +{ + uint ix = ref.offset >> uint(2); + Alloc param = a; + uint param_1 = ix + 0u; + uint raw0 = read_mem(param, param_1, v_202); + Alloc param_2 = a; + uint param_3 = ix + 1u; + uint raw1 = read_mem(param_2, param_3, v_202); + CmdFill s; + s.tile_ref = raw0; + s.backdrop = int(raw1); + return s; +} + +static inline __attribute__((always_inline)) +CmdFill Cmd_Fill_read(thread const Alloc& a, thread const CmdRef& ref, device Memory& v_202) +{ + Alloc param = a; + CmdFillRef param_1 = CmdFillRef{ ref.offset + 4u }; + return CmdFill_read(param, param_1, v_202); +} + +static inline __attribute__((always_inline)) +CmdAlpha CmdAlpha_read(thread const Alloc& a, thread const CmdAlphaRef& ref, device Memory& v_202) +{ + uint ix = ref.offset >> uint(2); + Alloc param = a; + uint param_1 = ix + 0u; + uint raw0 = read_mem(param, param_1, v_202); + CmdAlpha s; + s.alpha = as_type(raw0); + return s; +} + +static inline __attribute__((always_inline)) +CmdAlpha Cmd_Alpha_read(thread const Alloc& a, thread const CmdRef& ref, device Memory& v_202) +{ + Alloc param = a; + CmdAlphaRef param_1 = CmdAlphaRef{ ref.offset + 4u }; + return CmdAlpha_read(param, param_1, v_202); +} + +static inline __attribute__((always_inline)) +CmdColor CmdColor_read(thread const Alloc& a, thread const CmdColorRef& ref, device Memory& v_202) +{ + uint ix = ref.offset >> uint(2); + Alloc param = a; + uint param_1 = ix + 0u; + uint raw0 = read_mem(param, param_1, v_202); + CmdColor s; + s.rgba_color = raw0; + return s; +} + +static inline __attribute__((always_inline)) +CmdColor Cmd_Color_read(thread const Alloc& a, thread const CmdRef& ref, device Memory& v_202) +{ + Alloc param = a; + CmdColorRef param_1 = CmdColorRef{ ref.offset + 4u }; + return CmdColor_read(param, param_1, v_202); +} + +static inline __attribute__((always_inline)) +float3 fromsRGB(thread const float3& srgb) +{ + bool3 cutoff = srgb >= float3(0.040449999272823333740234375); + float3 below = srgb / float3(12.9200000762939453125); + float3 above = pow((srgb + float3(0.054999999701976776123046875)) / float3(1.05499994754791259765625), float3(2.400000095367431640625)); + return select(below, above, cutoff); +} + +static inline __attribute__((always_inline)) +float4 unpacksRGB(thread const uint& srgba) +{ + float4 color = unpack_unorm4x8_to_float(srgba).wzyx; + float3 param = color.xyz; + return float4(fromsRGB(param), color.w); +} + +static inline __attribute__((always_inline)) +CmdLinGrad CmdLinGrad_read(thread const Alloc& a, thread const CmdLinGradRef& ref, device Memory& v_202) +{ + uint ix = ref.offset >> uint(2); + Alloc param = a; + uint param_1 = ix + 0u; + uint raw0 = read_mem(param, param_1, v_202); + Alloc param_2 = a; + uint param_3 = ix + 1u; + uint raw1 = read_mem(param_2, param_3, v_202); + Alloc param_4 = a; + uint param_5 = ix + 2u; + uint raw2 = read_mem(param_4, param_5, v_202); + Alloc param_6 = a; + uint param_7 = ix + 3u; + uint raw3 = read_mem(param_6, param_7, v_202); + CmdLinGrad s; + s.index = raw0; + s.line_x = as_type(raw1); + s.line_y = as_type(raw2); + s.line_c = as_type(raw3); + return s; +} + +static inline __attribute__((always_inline)) +CmdLinGrad Cmd_LinGrad_read(thread const Alloc& a, thread const CmdRef& ref, device Memory& v_202) +{ + Alloc param = a; + CmdLinGradRef param_1 = CmdLinGradRef{ ref.offset + 4u }; + return CmdLinGrad_read(param, param_1, v_202); +} + +static inline __attribute__((always_inline)) +CmdImage CmdImage_read(thread const Alloc& a, thread const CmdImageRef& ref, device Memory& v_202) +{ + uint ix = ref.offset >> uint(2); + Alloc param = a; + uint param_1 = ix + 0u; + uint raw0 = read_mem(param, param_1, v_202); + Alloc param_2 = a; + uint param_3 = ix + 1u; + uint raw1 = read_mem(param_2, param_3, v_202); + CmdImage s; + s.index = raw0; + s.offset = int2(int(raw1 << uint(16)) >> 16, int(raw1) >> 16); + return s; +} + +static inline __attribute__((always_inline)) +CmdImage Cmd_Image_read(thread const Alloc& a, thread const CmdRef& ref, device Memory& v_202) +{ + Alloc param = a; + CmdImageRef param_1 = CmdImageRef{ ref.offset + 4u }; + return CmdImage_read(param, param_1, v_202); +} + +static inline __attribute__((always_inline)) +spvUnsafeArray fillImage(thread const uint2& xy, thread const CmdImage& cmd_img, thread texture2d image_atlas) +{ + spvUnsafeArray rgba; + for (uint i = 0u; i < 8u; i++) + { + uint param = i; + int2 uv = int2(xy + chunk_offset(param)) + cmd_img.offset; + float4 fg_rgba = image_atlas.read(uint2(uv)); + float3 param_1 = fg_rgba.xyz; + float3 _695 = fromsRGB(param_1); + fg_rgba.x = _695.x; + fg_rgba.y = _695.y; + fg_rgba.z = _695.z; + rgba[i] = fg_rgba; + } + return rgba; +} + +static inline __attribute__((always_inline)) +float3 tosRGB(thread const float3& rgb) +{ + bool3 cutoff = rgb >= float3(0.003130800090730190277099609375); + float3 below = float3(12.9200000762939453125) * rgb; + float3 above = (float3(1.05499994754791259765625) * pow(rgb, float3(0.416660010814666748046875))) - float3(0.054999999701976776123046875); + return select(below, above, cutoff); +} + +static inline __attribute__((always_inline)) +uint packsRGB(thread float4& rgba) +{ + float3 param = rgba.xyz; + rgba = float4(tosRGB(param), rgba.w); + return pack_float_to_unorm4x8(rgba.wzyx); +} + +static inline __attribute__((always_inline)) +CmdJump CmdJump_read(thread const Alloc& a, thread const CmdJumpRef& ref, device Memory& v_202) +{ + uint ix = ref.offset >> uint(2); + Alloc param = a; + uint param_1 = ix + 0u; + uint raw0 = read_mem(param, param_1, v_202); + CmdJump s; + s.new_ref = raw0; + return s; +} + +static inline __attribute__((always_inline)) +CmdJump Cmd_Jump_read(thread const Alloc& a, thread const CmdRef& ref, device Memory& v_202) +{ + Alloc param = a; + CmdJumpRef param_1 = CmdJumpRef{ ref.offset + 4u }; + return CmdJump_read(param, param_1, v_202); +} + +kernel void main0(device Memory& v_202 [[buffer(0)]], const device ConfigBuf& _723 [[buffer(1)]], texture2d image [[texture(2)]], texture2d image_atlas [[texture(3)]], texture2d gradients [[texture(4)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]]) +{ + uint tile_ix = (gl_WorkGroupID.y * _723.conf.width_in_tiles) + gl_WorkGroupID.x; + Alloc param; + param.offset = _723.conf.ptcl_alloc.offset; + uint param_1 = tile_ix * 1024u; + uint param_2 = 1024u; + Alloc cmd_alloc = slice_mem(param, param_1, param_2); + CmdRef cmd_ref = CmdRef{ cmd_alloc.offset }; + uint2 xy_uint = uint2(gl_LocalInvocationID.x + (16u * gl_WorkGroupID.x), gl_LocalInvocationID.y + (16u * gl_WorkGroupID.y)); + float2 xy = float2(xy_uint); + spvUnsafeArray rgba; + for (uint i = 0u; i < 8u; i++) + { + rgba[i] = float4(0.0); + } + uint clip_depth = 0u; + bool mem_ok = v_202.mem_error == 0u; + spvUnsafeArray df; + TileSegRef tile_seg_ref; + spvUnsafeArray area; + spvUnsafeArray, 128> blend_stack; + spvUnsafeArray, 128> blend_alpha_stack; + while (mem_ok) + { + Alloc param_3 = cmd_alloc; + CmdRef param_4 = cmd_ref; + uint tag = Cmd_tag(param_3, param_4, v_202).tag; + if (tag == 0u) + { + break; + } + switch (tag) + { + case 2u: + { + Alloc param_5 = cmd_alloc; + CmdRef param_6 = cmd_ref; + CmdStroke stroke = Cmd_Stroke_read(param_5, param_6, v_202); + for (uint k = 0u; k < 8u; k++) + { + df[k] = 1000000000.0; + } + tile_seg_ref = TileSegRef{ stroke.tile_ref }; + do + { + uint param_7 = tile_seg_ref.offset; + uint param_8 = 24u; + bool param_9 = mem_ok; + Alloc param_10 = new_alloc(param_7, param_8, param_9); + TileSegRef param_11 = tile_seg_ref; + TileSeg seg = TileSeg_read(param_10, param_11, v_202); + float2 line_vec = seg.vector; + for (uint k_1 = 0u; k_1 < 8u; k_1++) + { + float2 dpos = (xy + float2(0.5)) - seg.origin; + uint param_12 = k_1; + dpos += float2(chunk_offset(param_12)); + float t = fast::clamp(dot(line_vec, dpos) / dot(line_vec, line_vec), 0.0, 1.0); + df[k_1] = fast::min(df[k_1], length((line_vec * t) - dpos)); + } + tile_seg_ref = seg.next; + } while (tile_seg_ref.offset != 0u); + for (uint k_2 = 0u; k_2 < 8u; k_2++) + { + area[k_2] = fast::clamp((stroke.half_width + 0.5) - df[k_2], 0.0, 1.0); + } + cmd_ref.offset += 12u; + break; + } + case 1u: + { + Alloc param_13 = cmd_alloc; + CmdRef param_14 = cmd_ref; + CmdFill fill = Cmd_Fill_read(param_13, param_14, v_202); + for (uint k_3 = 0u; k_3 < 8u; k_3++) + { + area[k_3] = float(fill.backdrop); + } + tile_seg_ref = TileSegRef{ fill.tile_ref }; + do + { + uint param_15 = tile_seg_ref.offset; + uint param_16 = 24u; + bool param_17 = mem_ok; + Alloc param_18 = new_alloc(param_15, param_16, param_17); + TileSegRef param_19 = tile_seg_ref; + TileSeg seg_1 = TileSeg_read(param_18, param_19, v_202); + for (uint k_4 = 0u; k_4 < 8u; k_4++) + { + uint param_20 = k_4; + float2 my_xy = xy + float2(chunk_offset(param_20)); + float2 start = seg_1.origin - my_xy; + float2 end = start + seg_1.vector; + float2 window = fast::clamp(float2(start.y, end.y), float2(0.0), float2(1.0)); + if ((isunordered(window.x, window.y) || window.x != window.y)) + { + float2 t_1 = (window - float2(start.y)) / float2(seg_1.vector.y); + float2 xs = float2(mix(start.x, end.x, t_1.x), mix(start.x, end.x, t_1.y)); + float xmin = fast::min(fast::min(xs.x, xs.y), 1.0) - 9.9999999747524270787835121154785e-07; + float xmax = fast::max(xs.x, xs.y); + float b = fast::min(xmax, 1.0); + float c = fast::max(b, 0.0); + float d = fast::max(xmin, 0.0); + float a = ((b + (0.5 * ((d * d) - (c * c)))) - xmin) / (xmax - xmin); + area[k_4] += (a * (window.x - window.y)); + } + area[k_4] += (sign(seg_1.vector.x) * fast::clamp((my_xy.y - seg_1.y_edge) + 1.0, 0.0, 1.0)); + } + tile_seg_ref = seg_1.next; + } while (tile_seg_ref.offset != 0u); + for (uint k_5 = 0u; k_5 < 8u; k_5++) + { + area[k_5] = fast::min(abs(area[k_5]), 1.0); + } + cmd_ref.offset += 12u; + break; + } + case 3u: + { + for (uint k_6 = 0u; k_6 < 8u; k_6++) + { + area[k_6] = 1.0; + } + cmd_ref.offset += 4u; + break; + } + case 4u: + { + Alloc param_21 = cmd_alloc; + CmdRef param_22 = cmd_ref; + CmdAlpha alpha = Cmd_Alpha_read(param_21, param_22, v_202); + for (uint k_7 = 0u; k_7 < 8u; k_7++) + { + area[k_7] = alpha.alpha; + } + cmd_ref.offset += 8u; + break; + } + case 5u: + { + Alloc param_23 = cmd_alloc; + CmdRef param_24 = cmd_ref; + CmdColor color = Cmd_Color_read(param_23, param_24, v_202); + uint param_25 = color.rgba_color; + float4 fg = unpacksRGB(param_25); + for (uint k_8 = 0u; k_8 < 8u; k_8++) + { + float4 fg_k = fg * area[k_8]; + rgba[k_8] = (rgba[k_8] * (1.0 - fg_k.w)) + fg_k; + } + cmd_ref.offset += 8u; + break; + } + case 6u: + { + Alloc param_26 = cmd_alloc; + CmdRef param_27 = cmd_ref; + CmdLinGrad lin = Cmd_LinGrad_read(param_26, param_27, v_202); + float d_1 = ((lin.line_x * xy.x) + (lin.line_y * xy.y)) + lin.line_c; + for (uint k_9 = 0u; k_9 < 8u; k_9++) + { + uint param_28 = k_9; + float2 chunk_xy = float2(chunk_offset(param_28)); + float my_d = (d_1 + (lin.line_x * chunk_xy.x)) + (lin.line_y * chunk_xy.y); + int x = int(round(fast::clamp(my_d, 0.0, 1.0) * 511.0)); + float4 fg_rgba = gradients.read(uint2(int2(x, int(lin.index)))); + float3 param_29 = fg_rgba.xyz; + float3 _1298 = fromsRGB(param_29); + fg_rgba.x = _1298.x; + fg_rgba.y = _1298.y; + fg_rgba.z = _1298.z; + rgba[k_9] = fg_rgba; + } + cmd_ref.offset += 20u; + break; + } + case 7u: + { + Alloc param_30 = cmd_alloc; + CmdRef param_31 = cmd_ref; + CmdImage fill_img = Cmd_Image_read(param_30, param_31, v_202); + uint2 param_32 = xy_uint; + CmdImage param_33 = fill_img; + spvUnsafeArray img; + img = fillImage(param_32, param_33, image_atlas); + for (uint k_10 = 0u; k_10 < 8u; k_10++) + { + float4 fg_k_1 = img[k_10] * area[k_10]; + rgba[k_10] = (rgba[k_10] * (1.0 - fg_k_1.w)) + fg_k_1; + } + cmd_ref.offset += 12u; + break; + } + case 8u: + { + for (uint k_11 = 0u; k_11 < 8u; k_11++) + { + uint d_2 = min(clip_depth, 127u); + float4 param_34 = float4(rgba[k_11]); + uint _1390 = packsRGB(param_34); + blend_stack[d_2][k_11] = _1390; + blend_alpha_stack[d_2][k_11] = fast::clamp(abs(area[k_11]), 0.0, 1.0); + rgba[k_11] = float4(0.0); + } + clip_depth++; + cmd_ref.offset += 4u; + break; + } + case 9u: + { + clip_depth--; + for (uint k_12 = 0u; k_12 < 8u; k_12++) + { + uint d_3 = min(clip_depth, 127u); + uint param_35 = blend_stack[d_3][k_12]; + float4 bg = unpacksRGB(param_35); + float4 fg_1 = (rgba[k_12] * area[k_12]) * blend_alpha_stack[d_3][k_12]; + rgba[k_12] = (bg * (1.0 - fg_1.w)) + fg_1; + } + cmd_ref.offset += 4u; + break; + } + case 10u: + { + Alloc param_36 = cmd_alloc; + CmdRef param_37 = cmd_ref; + cmd_ref = CmdRef{ Cmd_Jump_read(param_36, param_37, v_202).new_ref }; + cmd_alloc.offset = cmd_ref.offset; + break; + } + } + } + for (uint i_1 = 0u; i_1 < 8u; i_1++) + { + uint param_38 = i_1; + image.write(float4(rgba[i_1].w), uint2(int2(xy_uint + chunk_offset(param_38)))); + } +} + diff --git a/piet-gpu/shader/gen/kernel4_gray.spv b/piet-gpu/shader/gen/kernel4_gray.spv new file mode 100644 index 0000000000000000000000000000000000000000..61e5b1cc0b7161e3317f50b0f55fdf0ac03c307e GIT binary patch literal 38880 zcmb822b^71y|y>ZB!SSSBOoO}kluRMrN(GofSg^+#3JV%6wsiRuR< zs`^-=S{$XSthb_Gsv24I2?rg1(B^Xodp6&8yDfEGw(78|ZOc`oD#JS|v%35G)t_Iy zs$QooPrUjHlqD!)^LOxnn5(BDo%CSBL0uCjPMmPizFkvh&+eK#c-qX~u9o<1ZR=e=O$y@D$t+QHsm~Usb418|i>Ga0fPUe9z_s;4Xm?0lstqdO=m^Wp5 z7wXu#gLCFh8QiSyz-IHdsx!CZR>ZGvugS9#^`btFuGWCh>Fw_MH)-njS8L0&`mj7p zRqMd#&g`4g`~OYEVDTBjdaOm=UE|JbZR$n+>8#!XUo`HK)q3!~XZ0LDIA>r+?_s@D zq4oT#dn}3_QEdQU=nEdn|EYNl_OZ}&65rC@AJe;MPVG9SuV-+&eOsa07;Rw=0=j$t z_vALSz0i$r)st>Jj&CFC*7!QAv0$4W!T)UnCTM)$%$fB_+$U-ti^kDeZ4Phz>M_VY z+Q*>o7$$d5nb9+6V75LZI1{#}ZsPyr2%63e>%S*n`tP}!{%_mLyZ6A%fjJJq{;PQ` z>i>vpdpPaae%f3IHIF&dCU;9^Hl>OZuP9GZI*8WoS=K@ds=;p*e^AYa-0+3^r80Crk|Qe zIZidlxhQt|>Ogp6|6}jdlcMkY6Q8~x08ZZzZsi>}bN2LZSIIu9c`RywL^Tmk`?a4o z{nR|V4W@SQA5I;=|Jcv8TJ-Z!^VPXJs>7($&m&uT512P=w%x4TSMylZ&(7*-IPL3K z`^Y`K?zeI+UmXu`eEwq>n+BU6I>wIbSmM*iFGiUJi+Dyr&VVf0=V$k#qjSoXb~dG&!3K#($BuHK(T+dEc`z?{d|GCU1Qz{})-E)3b}5=M*_F zY;wA%|BICMoSt9gykKF@5!J=uBe=I8-aBpKzE|@o_x)w?#;@Mza*u&IebaoBbnmMz zMw!opp!cU}Uf)ma>g}1v(^xdG@BO`}@?_^cE?-?m-Q@rGLtE6@VdwM}#AnZ332vP) zBdY7cg9CF9o3Pi|`BS&?S+S$KPW!aU;E~l0;Hh&4W-aotKb_V4(dX*Kq2I04^ZIA= zq)>a-4`ce>()hL79$9?|zR+v_)NNbOY@OBZ=-#`j$TzCGn|jLhdHplGTC?oqY|GYB z-392Y{g$pCp`OaQUhku^^XF~7%c7^JT4NSAo8Py!(T?pQ{7>aRYTq6PPvN@LH)~pL z`zXBC&(7*I;Mv`Cx;^*xePS5zsJ;OAJkjVtb{4a}+U&m(J&)Owz7p3i+o(QmOv z@0~NpEl{twx_f{p?K+43d=Bm*{vLH7&jYsgBk-K=Q+$M~`A?|Je*UbD{{oJ`k6E?< zt8Mn*4C5WuZ{e-$a!2(yaN7MBxNO&fl~xt^V~2juCuyn81JZVg}3sItUe6pnW3w@ zzo%=_^F(KL_pq4G>fSc~cpLvz8-Kcu54G_xxACvE@o%;9@3iqBxAC90@n5&`*V_1> z+W23F@s8^6@XYU<;BtN!<3dxOw;gSK1RVc``(ueV`{*{lbjyF?zF)4*zCs&cvE{!_ z_Wdet_SM?>8ZH0&W1!c+wc70OXyfa){I^T~4chD*4dWfvJK?SU)LHEfo-%Kc8&2yS z-VZ*xcjkc4RknRFynFJ1U*XjJgf{;vZT?f+{O9W5`e<6mp8_w};=EzJqq-U1y3ch~ z9|q6khGReO9Ol<83id!WZ3dxrf0A{SJeT z9! z-dVl2rI1<{rttU)_-UYm=_N{w_OK$Fyud8%^)+n>KxLkumu? zvDM>fzLbhS0(~;6@?BN*k?6Arr_78Y?T=D#^(v-?eO>}vcYpst49T@5`rw@I{<)1I zy5H4Jo;+~sBL5}oF)oGP!^t*lpnssRC)U<8fpg$@b`0K(v^}DZS*F(P*`m2x7OmB8 zRzX{;&Er(5`=;hGg>|>~ zx#b&;X3Y1%K5J$E-AgZ{mHYn}ZTy#Q{8vqm*0XD8^&0%t)4JyM^$*g*67~1A?Ae0` zrgYETzkfc}VBbK$JGa%Jf7bpj{1&)%Zgp0p*iehkg%Q=V@F_F&;^~f;ZQGE=aIeHeH@jo}r zud{lgjen($f31yweHibkUV>+=KLpP@t;>%O^;lmX=G$5Q96oojdrrptD!jM9r}q6V z_!Pbb^Jd~)ya|ur5?r)SWd>cp)x@(}?;hmq*1X%V2WpnWxhi;?C+aeHMPN)mdE*FX#KpVZ5Wdx`@34T;9j-9OldO*D&5u-3On_#oh7T z-)4Ilp80*O_Un??eg7()y;sM*I4rKS`uZ@=I>CFIxSxYtIY(5lfm@#=>t(BRy^dCn z^$j@d+LLzQf;-{%cO*BC^4>ZMK8#Q~##yhGl;hD1?!Srw8oE^5%)=?b?c2Hh_Cxa)?WH<6oiVx{^ zzh|K>>i4K>&O)DhXEdIzwYV5ttIdw;z2H{fj_O)tnzel&cu~HN>Lz&B{jp)T&gzrH zct`aNJlEqF#6G_?bN6E5JApmzcNQMv>sqkJeplh4?)MWOV!xa45Z~T{HTJs$j|r^v z7pX^5So&%|X2;Lu9_zQNg__SRYNa2yl3Me*z&4EK^se*zIetBghwaGMty#-=AA>oJ zn^0h%J)<93;DZY6^Jsh~75MN1A5q|=3w%t0k1g==1@0>Fk!f0?#h6&(g{7b9C^j1wO67rx*B)0-ssnvkH86fzK(h&**9E`~qK4 z;0p_Uae*%>@TCR5tiYES_=*BwS>US*?6ZHXKinh1>$PaT*W~WC)?VzW?JL6FYuZ+8 zu%8V))E_PICkp&nfuAh!=L-D!g}Af+{IC=`*PwVl<#R%LKE2rZw`=X^n0~tI=Q(CR zHOHmqb8*_-bPd&GD;k zT4?rF%`foMuFv#p+ZLM7^lCd6n$PlTyB3<`RNJ%Ad~R3Ur`9^_=XGG^+V96TD9`E{ zd$=(k?gO=s8h6&u@$1zVZ7uG;a-HO_jayTB(G*yoJIK2%_zE22MM;7=F$GX;LCz)u(WnF0?L`11w+LV7=mYZTaLio|YFV4oSH?^5913%o~x_b%|h z1s-4E2?gG-!21_?T7iA`NFRLe2tK92rxy6M0{hGn{~HVJ^G5Vr3w&FFZ!hqj1-`4m z_Z0ZP0^eWY2MYXPfgf6k`RoBM*Ie$JbGp>)DbJdVS8w!hgFQUAz7O`?(zYUz#;(g= z@RSj+jE~o{btoR|z`c%H)K|jGYnZy%Ep_?YWbm4$&$|47+mI(#Ut_$MsA*f5>+ajA zjgPLjD%?Daw2_t9P-CK*!#16dR)0I|Ym~a$D%i$QuR~dzb~m8bPt7(qq*gQE`V@0* zLY-V*W6h=JoNPwD4rMWl`8KE4SKWMDP^+2GSo3-PO+LRsG@qI_e_x`_*d3_l#_Fft zYjR?D0~@QhHmbkt(B@cD%(Dkvn|X{e&py=2voF{@YO6Aa@znO;b$3h?sP$Dh-+t6; z<}=oO2U92CAz<^VX`4uGe~g_(EjLy_?MG54_9(EiYTAx&Y{nh~mK&>|_IFb!_5`r8 zYWCaT0BExqe5N!Z0{f;(>mi39|XZGDGWgBO~GhgvL4=rO+ zb1WBAXDpY19gCW4aw)ZQ*jkg3^*dg*==wRH%c(D;*uFW8y_(us&t=cGYpBhwzApUQ z2CM6T6Se-G6zw-ts~Kbb2OF%W{Wfa*t#14s4Y#k_@1#~s{JjlUGv9p;pFq6{++6ol zmY}Q&e}LMy)Ewi3)N1B2))@Ee$0!~fGg4zeNlCsT4IFD>`y#b%tWB;jHCV20f1ahb z@7i{0_;Zbxww`ae{x3Ax9BJ#arnQFv!zU!Ts6XyP{pJ3jZR`L=Ahd)fUn4(|G;|9drdb=z~E_64uG`N6wv%1XMA z$HUc9;2rjYANwM678W0 z-Mhu_D7acm{7yvsASc5OVrT#IMn?(3|@b8z#k8!z`Aj+*gbZm=<4)BHTv!YkFErNm{? z{_jfZTViwj*}>0ew)=Hz{#R@EOYVCCHRsTC=lc+I$gO{8>65MC=*!$py zY|2^Wd2$k5%@Ut8YoBxWNRGSVOXC~+I~eO{cerr7@WnkL&v|R_>;pd!cqD#p_(}zT z7yN9vV>$us96F}>ccGor*iQuOqvn{qsU5Sr@w2HNf9xkWc6Imi^VH6n^Isk;LH+x4 zbAA~g4}Bej+~->l&lC5k&$af^>(MLJ+cSUWU@GdvDf=H3!^~kHzl4Led+_LLmUiD8g)fP}?~KB&ePxc6q?A%)vN z-ywxNU%o>MxAq-UxbyLB!PR}26uY(WlER%2-zA0Hp6`;voe$q7gfXJg9>+keFqh8?R%(j*UxuR;nu!~3OC;OP~kpX`W`CW{`q}U zxSt99ZYbR6ZQnbE+h5-~h1=ha3hwpG?}uW)7rXDB!p*l?!PU1cxc%{ar1)F=?y2Oy ze+u_L<2$Hu<7XAz`;6ZM#cu68r;@+7;KutsP-*x5Q@D1&4+?iadzVS29^M&6_Xxh#vmvP?$8>^oD z-v*n1sEGRxT>YFT?z>=P)pxIW*F>$B{`>%}mUeywR`Z$4KKeZM6H5B%vzM_xU;VVP z9on?@vqqykbswFZpQCx7b$qrj*H7&I*7n5SYi)BUW^^>ZuYlK~Hda6F`gsr6<{X&+ zRj``*Gp1j|)hE$UzeD&9Sj~9v;quJkZ^1qvmUH+!xO#lv0JqQcAJEh@Z+`^;o}!+) z{}b3;+LG(fVB0JA#9!d*@%by*_DY}sfvcyzzkzK}J?;G+Y%Xn+n4dcs`#-?;L!0CG znMKXH(*7^7+HAPbqi=%M^!0g0F7{bw2>aiB>f*lPvzGN=h+CYxgW{pRv$0!eO-6vt z@d*9$GvP?Mn$L89Cx?77in+`!PV5q3W1Sz*)1}~}DC+i2?lZe>+YjTd6TdXr_{`HX zaJ4tcm9eRnW3wOHd`5HJw(a@4Jp4IoW0QXcH1+$MvAqqfW}G>Eu2VDriq!Vc+PPbq zdKHR?_Ej6ZwfUW^)xhR(p2RCrj5Ci|KhK%BgFP2)(`QAwHqU{z!0vr*J~zsD;{5Ky zwl;VQwK3XqKEDI3&BHjKN7XX!b--#F_qt#;$Ng9Q*9UvpZ*A*Q)a}VubZ1td!5*n+PU0}+GnxN zshzVesO^vWGLPGUGml$?<(bFrz?ny%r{&J$2ISfvJd4^GZO++N)aqiNx1Afuy(8Fm zt=$JZQ}06Y(7tP9x6Zum1~x~|q20Anat_INqL|Cv;>7L+Hr9RYv%@}Oin@K1?@h67 z`(eCw;`ap`pL1wDSS{y}HnnnW_QUp^|NX$uMR}dxAFl3o`eJ+z0IS&_KYPf{cObR> zv39Nwrapw?q5Yu7Zk_fH1Dm64?{K)f`46R*i;tk5L~*R{k)yz#Z`zKemgl~7v^IQ< zH%_~L$53lChhsSoZ2Q`drIx!-F2i;__$q2+wCOXETHXHp?DuZ4TCOu*n@@nN>;Eol zx!BK1*;glmohxhSYBKc{iidO6)7Y)Ew!L6;WM54Mt7TuwyD8=}x7b+wG#%`HVgdb@ z&w#6)M9ElYf*niNLcdvX_4Aszez39Xj%gaTTKYH}Y#+4^P|J5B_6~BL48EJX)|-CK zK~qmZ=Yq|dehz}w^hAs@8-E>bj-jnn!TPAj=QMB`b9%!Q^B%Z9>K8Y8 z&H$TB+tTblbDstF+O`Qr`;wG1DfV+7wRYR|`gS(hc?mzK;m*D1`2x7-xw<~)KNqZ? zexC^SzKXkVHuf%+wl?lsfC z=$HOpDyBFt$1B(0z3V=9|5<0emw_FBu1lAL)pA|Zre^$!)W*AyZO_l#SAgA{evfWW z&)qBG>b5_DT5kLy)W%yI?{)1eus+$_uE~4h>i6U0_2g=>+7Q@#q5MF~!4&gZ+txMI z`qrg`U1u!{{UF6kCHqe1ebYkgS(#U`WSaRSlzuN{}9+*>Ny8K49+>AJ?B9BcL&-{ z6#MPm$&GPtoF~U;oqpX3wqMyNcY)QiPqe9(bMHK5o$mpA-!IqsUNrUOybr8)HzjfR zgA->z<@(vqmelrhBWn9Onz?uwoM(@Rz;bOq!+ivtXOBm~a_=d=U-~F`I<+y{Z0kX4 zb@9ijhp^l3$HBI1?cVwX^&)LL!RE-``V?3#drST(#a!kVC-&1|W1SPf zr}zxmzGZK1M!TPds}Im0_rw!mwXFA(VAtDRjz_MK{k8A5X`MEo0^4TR=5t`RthY9` zaz5;P*6nGqeJt1e8Mu0$SBAi9_Qz|0+wm;U+?-!}RMDfu6g~o23_MQiuqipZX zaCP%PM=ck>Ks|}#SiQb{6}((y`wF!@&nqu#!^e2zw438=)Y{D9HSZf>+t>DWYI&Yl zz6thu#Tae+JWH)Eeu+AJ>sw&QZtd8=L;YQfhhzU($ z++t(BH~k3g^G9<17;Y~6^)j_Q&j~*P`<$Sj&$~Ybt1n=0n(t>|bE=!;htz7Wlh0+p z06X@qlQF-9tEW$|fQ`$&@mFAd)Xn*GYPIw3@x;bB^ zRx^k9{ojGj`3l%}Jg0sSSC7vdVB<=kKfu)|HS_UDuyN|q{sdmKDbfB6KeH)b!|n55 z;JJpI$2k4$?{BHKr!Rj6tGSQT-~WND$LDX2PubtUqp3eee?1re0X9zE&yWA2_VB)| z?Vl7i_mMbj@D|wi(-!XkYyK8PQ_q+e2di1K)*WDTY4iIt*P;`g^(^~70$p43jRdP% zd_HszM}f_$&3Sv1+-m9Pl3>?)2;a=xXf*ZsECqI*OP{6D)Z?=Z*xaShvS{kLCoczf zJ-y%9hB@>}j^)AaeOUobJw9&(xA$d5H1+tb1hy~bT&;|zp1!OCPG4-p9QtG~R|VT& z{hT-7ZK$P<)xl~*#k{QnSC7xz!TOZ_SQD%g`9 zF4z9~u0<_vuLo}L-}-3k@fp+jl>OTPO+9_s5N!X(QjFJbd+Snb%bpnvR!iK*;EYR~ zeK*c`KH3tu30N(0n}XAxHm?)L`EE#C;x-36=EQ9QPMo&H`3^~2;L>!%){-N4R&`0n5`W)E!?^?Y{R6KpPR&av;;)Xd@Mka1w|g~_=WT%YWpy}|9* z+-?=G^AF`UJRIo~yN~8J}lCpWAHD&)Z$t zoTu{h_K9ffw(m2j+<2dvZQt7XQN(v+(#}Z`NxbIG=x$b3U4Sa-ITKySwS9>wGFaarRTLpY3dkpZy$5t?hL5 zA!4+>2dvH7bv%>$EQ*Krvm3j0`g0E099hQ&V708H{0xe@%q>prxnN^6zvqFSUv>K? zcm8eLei(0^`18TWXD%)Pt7U$*sg+~1ADQ0^!I`Jzya-J_IWGpQ`P-?)T>?*>{gmrx zJDJ~&>9;vYGj~^D%l8tOgXP-%Ug9dSxy*AVSng-XiP+u?o!VlH!w6MG}r*sR@6VEdM9%NW|d z8Ll3mTjb1(vG!lCkNvZ+wqcz%ZUx&$*5U(TwXCl;wQ~ISHS6?2aMmd~Z$nd0&fCFi zjyZ83f+x;?%Js9ItnWLE^}P#Q*7r`Z-1c1Gd%#)WyTNkT*WWGQ3qFO~7;Uz72eo?o z@&H&Z^XfJGLAbj9_fyOLobK;09|GS>t*-q(YBll0)I-=D$0K0p+uHr|QRE9tE2t`{fg0wd@!9M=0hpx7b+w^*GqRB-baw=E~mq6xh9?uFn`^J`Gm)yL5AW z25cU6$MYDqT5^6CY|ii}8g3iD`+gEWiZ;~sG3QfY_4MO&VDs+7`JR3}4Oh>z$TMKC zOX}&{5ZGMWlKb;ub1i84{{^sq>gm%L!S>0%J0`h)u7T^}JX&W=UjjSE?B8d>YB?vg zsg-N&df1-#hUdW9+sXMnntF168LU>G6EDCMXFuio*-p-hO&Etcz2JyxnBkwr@nx>A49$$gKbk=a{UBsdwHfC zgWpf#>WTXq*jRPv<0Wb}@z1G;h;bc$0d~JwdrrSX{VR%x_E#Ibb=LmZU~}Z0{tZ|y z=d}En6mywdoY>!jjkO=IQ_J&x;qSnA)2@2HFZ?}NUH{jpi|E4p)!QKftzEuKPdX>S^y^ zVB^%&-kV_C)aKaVpjJzt*&el)&oGO@^BG24`)8OA?G%5zp6A(4HRi#hy`S~>-M|R2 zHrsuRHs!Ya5c_utu)lYg>$@aaAIIypVl>=2DepZ?p{bwI#4HU~n?%fbVwM4$N8Nwb z(rfIpVB6AWKm2W_n(b;|9;`NzPXbrN{d-$Ai-%*PyLGI;!_@9R)7QVpy@6Ui>$D=+ zd_(B|9%l^oN^pHW?1MgPiCGzJjBA^iRp7?hzA^f!C1zEyG5+0AVpfA2lm6t`)+vDG9KSFJc zHhuh^uDWLKioX)CaEZ|${p8|rN-9@@8S?ADpb?ZM{AwRH!uTCT0~ttsX*w>Yso zfsM_+*coi!)b;Va*#)egm|ek%@w@Zg;678S>oW#zcQGZ;Eqj2?qwbjfov~VS?g=($ z__&6rUwgsbuj=~Puf4(Q`q+{VzqZ5BvlkBsI|gg_;*r!xQ9O)2y0KfQKgWR0 zk-c~;8yKd_GxW3cC>bbs72djBR>uWUo6z!x&^Ew=DMx)UL_p8?5tIb3+r#Xz#$Fa=fKqxU)$h0hZcZy4kgFAXzIyv z9$2mX4(EJ$;vAP;zw+GMoPH$dh1hKS+-5&q1XnBjrDm>+sfSntb6o;Ys^gY9R|)ob86SGDC_eJ@yB;;sX$g5slF5*V8-$kfvzm{4pu{VQTZM2?!Zh^bk)b;WFzZI;W zm=AyxdEsVa5>%&!!us<-T|LO(N9})-wD>1+;@S^ zukILcqE<^gcZ1CKeh_S~1y%E1 z5LiF;oX-z~on!lDf8`HR+>7o@_mg${{|MOr=Un;-SS|Zrn_9Wo-IqB(KMKxymYg3$ zQ%}y1gVoCO^HF%>?5A8m+sVCg3&vqizi)UPTb?f;1IxAfeEBJ`xy-}fI+C2Qtl6gK0S2vIO!RE*{<~guht}*h@Q_N*< zv9Zrne}(cw!|$MtuY!I4PVTS4&22wlq?YGcN3yEpv}na5I;9heiJ9sHa0;AI-@-&QP3u^;}8eg5|PqhMpwuaAR` z&%SyTtWWlpTp#^y{}UALZ<1G>zWLs5Ws1k!DE4Os>J=Mon=4W5&nnc$`|pV)-t%M) zipQ#y#QQh<8Jqu>lJWj~Wf{+Bz~;{1nm^ITztHgX|Fdv&x}S`dXD*%u+lDst$UV0m z+fx+BXNOQJ$jwm6&I!J$%=x{xn6+K8qdmP{Y;J?w7%~>wbKJ`YV*@D9=;0 z#qXsKM(~kEPhJcT#7qzXvYY{znCO z@BO6U`u(hp|Dxdfzglqpe^c<+!M|;I=H>fv`&G`%58&!gbH0wG{vk!pwRGL=vs&iu zWw2W2?Z;p>_lZ8v+fOOZh4GF_E_OUWrnpCK>*o~DA$_#z!w^Ea$s}hd*|iBYR3D!__tAv zf1`<)>tnq8XvN0v9usGeZI5kdipLIhY25>NZ18r}J5k(YyHID3`QObc_s|*zchC6W z&MEEww{yb%_Z8M_c*egf+;+-6vl^Ou_RQ*Fwd@)Dpq4p$J6J7yW=*hKxo6gfJ2vOo zKFY=RPyejpI$+mOA8qp#BW`upD>?QIJ#+uIJV_6l+Oc-`6oz9+Tu&beId zeCwa~b^_a;KHBtg&D7K0F5t4g-Qa4zlSzAffYY99CO4PsrGMHR2ev(ZwCUrTtEaua z!DV~r2daN2Xt<>uOtTK}}SKiKy4(WZ}kMm_Bv1a^+ICk_U?PjZgQ^@;x> zVCN|J;)!52qsaV@~{GV70^_4mLjfO|DPkj|4l0#2*D#OZ?Gb z<9*+4`*MAZ_d0M4Si9E)upLP8*pHHH#QqK5m-+yT*NB6tbB#Ej9686Vy9(|( zIl184d)oNag6rQ`aQ$Z#{66@shG)Ir1-D=2Is9%k^_;^efYmam_E{}+dLmdY=WsVz ztvrXPz?}={)G^7$jz|COxn8h)P9JUhI4A09ZyLC4Z#rD<72@=9kDmljd(Mg6T+WC7 zX>TUj_Vm%Fk8`S?_WHs0A#*kWb{@0G<@&^bHrO0P&GY&ku-eI#JRi&jt3QK}@3sfQ zYWgJ4Jg|8ZGasy$m{Y)N#!Mp5c=DVIHkUU0dm8oW6#JWVO0J*nOd#exjom#h&Yqr# z?Qn|6Bue)5p$$HS`Y?)n`UvXm>9fe;-U~ms;O@~23a;Np1-}G*X~Q$$XTxo$+@t59 zsb`Na0IOw>+6T4F<#}MW?9ubVYULii5boIA1NKoawtxC(&s+?4&*-B~AIGSk_AUdf zCU$at+*e$~CzbO+9OJ9at@EVjt8pNAClxWli1>Rx8)! zMz~{hj_spdZ2$Dnn%oR_P4v;Gk7HC%dmjL+Wegt#JFcvOTp#_N=i3^)Yaq@VbYYu9 z@i>u^HRx{e3DlD*u0aoV*5JeB$Qn53cNbj!-iD`tcfie4&hedS>Y3xaz-rF1{r1{- z54@Y&Ij{}6*!=ou&h7&{XZmQ<$3CiO&h7`ReUW|Z_W=)r)gGY4=OM5@>C?ktwOp5+ zuSej^)Xt~#_c%E+pMKW2Q_nb`0;^@5>*4=7xO&F<4A{On zHrteoZBJjv=@>o_-k4e+ZH`kd*m0grah!9hGtMuQBjfbm@|A{X zoL_<)SB~>pH1&-0Ij~yBc`*LZ!__m+7r^$#vDv0vYWuSS8Rxe6zXVs$ zIKK_HFOJPNJ;RLeMb$94wA<1|Xfd3u8#=X)rQ^GxcD^T*`K zIQ@+L(}ri9KY$xoj`N3T>KW&cz-k%iuK2$USI;l5w8hV8?k5#c`fXopJt}92uw2`L8uRFYQh!|UJ!srAw3IMp&v?`;=SJkFzJoaZ;#ab7@i zoEK4NoPQ!m#(5O_UmBiq{vK{zInFoG)U%#{0IOx3N8sg8hi=$l@!PMUTVkb`2J3gjMMtff_vXtZ1Lgt zjy67`;QB97aPy6B<4YG@|K$p<{|aq<#e(a!H8<#9W%l6iUt9jmMKG%k)J@=&CT<%Bx z)80B@+tWvzKJIDtw6`9(Y;O!)ZT-f_`P>kk_T1BQbGe`OPkUp*wx^FaeLN@B)7~av z_fS6HZ35~!cL3Yw8%>*X{gQJ>uydW9JAu`V_jBydVB>Qh%k@e8E@0=* zXX}gc-xaK8yw~5|z{VfJTsmgCKE`|f-o3GVZ5QX-el@o1DIV8SytZFMeI3Q?`TMB7 zp6^A>-W0D}a((t|;|Dj~HqE;a+&oXv$8q2Z6t(db`)EG3_#X(?-@mai{vh}P6m|V= zPc8nF!20_eIQQuOinR{9?eRzHeje*3SM{KtT8GyGVvK61x% z9K|u*%hl0*$2XdO#yZz(nUw~YQxjFF1UFz*OTFD z-4w@WKDGGwfc5v^NwDv|a5d*!f7?@w|1_}v3mX6Fa5d*HIr_lr&U4meCOB)NeFnw+ z_Fr4_&jQ3vCg$x=K3c3dMm}_W{UgT`MibVTz`Ptx&9#a zT;j^PKBeKd?R{?$Zk)f%HSavQW7?CVZ6g0#s``AeHcJPbs`jI*$80{Wy;}=RNarK3u=d#|3aTi$^&h z#%j-eTnP3$rk?q@2yDJ`J}yR6&p0mu8|xA6(nhnrXqPpbu9=U^A;#;k-Q2E~ZEDN7 zt^liLpIix7v$XFM_n32d6*k-7fug_jtd@Dc1O0A_$DNey@4LXx^F7qrzt<3_ehKHU zxvzz*`HYl*EBiXQ`u#=R^>E`ZY~tPrS1aS*4_9w}4r%?)?*_Q}wApWY`uza;9-(+V zNU`7dQa{w-`=}qLIEIff;4}0{mpdT`zO{ zJnZlLZUNi2wZF0a0QCnc9@=kf?ADpz+rj3@Z!A9qR&y@>y_x)0in+`!Hr6$~1MJw7 z>n^yt?3ZJb=QovigI#lVpXZIg2W$>)K3j}OyBDnPdEi{#2X;>1pxBN+_oJzsNB#iV zICa~0e%17Ui25UxM;fmGN5T3`Cf;k=$H4li8}Iz98L$0Ouv*UTPk_}d9`1Fav96W& zoZF9qz1FGc+tP;!)DrV4uraS-Pt2#`#@N0w`lu!5Ghk!f^NIN^ z+!+2>$LOP$HGB%}8fNUD1DnH~?#-vc?mzu(Q?8Hu%f03vu};2cz~;+mxFN8b>uJ0; WwQ~Qu*Npdk`#jk8JQu~Di~kSi%`^i5 literal 0 HcmV?d00001