From 0cf370f9c704f04cfd9aae02b5c7b194f427ba84 Mon Sep 17 00:00:00 2001 From: Raph Levien Date: Wed, 19 Jan 2022 11:58:01 -0800 Subject: [PATCH] Mostly working rendering This exposes interfaces to render glyphs into a texture atlas. The main changes are: * Methods to plumb raw Metal GPU resources (device, texture, etc) into piet-gpu-hal objects. * A new glyph_render API specialized to rendering glyphs. This is basically the same as just painting to a canvas, but will allow better caching (and has more direct access to fonts, bypassing the Piet font type which is underdeveloped). * Ability to render to A8 target in addition to RGBA. WIP, there are some rough edges, not least of which is that the image format changes are only on mac and cause compile errors elsewhere. --- piet-gpu-hal/src/backend.rs | 11 +++- piet-gpu-hal/src/hub.rs | 11 ++-- piet-gpu-hal/src/lib.rs | 9 ++++ piet-gpu-hal/src/metal.rs | 15 ++++-- piet-gpu-hal/src/mux.rs | 14 +++-- piet-gpu/shader/build.ninja | 6 +++ piet-gpu/shader/gen/backdrop.dxil | Bin 4672 -> 4680 bytes piet-gpu/shader/gen/backdrop_lg.dxil | Bin 4676 -> 4684 bytes piet-gpu/shader/gen/bbox_clear.dxil | Bin 3160 -> 3168 bytes piet-gpu/shader/gen/binning.dxil | Bin 5800 -> 5800 bytes piet-gpu/shader/gen/coarse.dxil | Bin 10984 -> 10980 bytes piet-gpu/shader/gen/draw_leaf.dxil | Bin 6880 -> 6880 bytes piet-gpu/shader/gen/draw_reduce.dxil | Bin 3940 -> 3948 bytes piet-gpu/shader/gen/draw_root.dxil | Bin 3944 -> 3948 bytes piet-gpu/shader/gen/kernel4.dxil | Bin 10004 -> 10008 bytes piet-gpu/shader/gen/path_coarse.dxil | Bin 7064 -> 7068 bytes piet-gpu/shader/gen/pathseg.dxil | Bin 9592 -> 9596 bytes piet-gpu/shader/gen/pathtag_reduce.dxil | Bin 4644 -> 4640 bytes piet-gpu/shader/gen/pathtag_root.dxil | Bin 4716 -> 4720 bytes piet-gpu/shader/gen/tile_alloc.dxil | Bin 5048 -> 5056 bytes piet-gpu/shader/gen/transform_leaf.dxil | Bin 5664 -> 5664 bytes piet-gpu/shader/gen/transform_reduce.dxil | Bin 4696 -> 4696 bytes piet-gpu/shader/gen/transform_root.dxil | Bin 4824 -> 4828 bytes piet-gpu/shader/kernel4.comp | 10 ++++ piet-gpu/src/glyph_render.rs | 42 +++++++++++---- piet-gpu/src/lib.rs | 61 ++++++++++++++++++++-- 26 files changed, 153 insertions(+), 26 deletions(-) diff --git a/piet-gpu-hal/src/backend.rs b/piet-gpu-hal/src/backend.rs index a4422b9..1086d3b 100644 --- a/piet-gpu-hal/src/backend.rs +++ b/piet-gpu-hal/src/backend.rs @@ -16,7 +16,9 @@ //! The generic trait for backends to implement. -use crate::{BindType, BufferUsage, Error, GpuInfo, ImageLayout, MapMode, SamplerParams}; +use crate::{ + BindType, BufferUsage, Error, GpuInfo, ImageFormat, ImageLayout, MapMode, SamplerParams, +}; pub trait Device: Sized { type Buffer: 'static; @@ -47,7 +49,12 @@ pub trait Device: Sized { /// Maybe doesn't need result return? unsafe fn destroy_buffer(&self, buffer: &Self::Buffer) -> Result<(), Error>; - unsafe fn create_image2d(&self, width: u32, height: u32) -> Result; + unsafe fn create_image2d( + &self, + width: u32, + height: u32, + format: ImageFormat, + ) -> Result; /// Destroy an image. /// diff --git a/piet-gpu-hal/src/hub.rs b/piet-gpu-hal/src/hub.rs index ec4d169..ecfc2d4 100644 --- a/piet-gpu-hal/src/hub.rs +++ b/piet-gpu-hal/src/hub.rs @@ -13,7 +13,7 @@ use std::sync::{Arc, Mutex, Weak}; use bytemuck::Pod; use smallvec::SmallVec; -use crate::{mux, BackendType, BufWrite, MapMode}; +use crate::{mux, BackendType, BufWrite, ImageFormat, MapMode}; use crate::{BindType, BufferUsage, Error, GpuInfo, ImageLayout, SamplerParams}; @@ -312,8 +312,13 @@ impl Session { /// /// Currently this creates only a 2D image in RGBA8 format, with usage /// so that it can be accessed by shaders and used for transfer. - pub unsafe fn create_image2d(&self, width: u32, height: u32) -> Result { - let image = self.0.device.create_image2d(width, height)?; + pub unsafe fn create_image2d( + &self, + width: u32, + height: u32, + format: ImageFormat, + ) -> Result { + let image = self.0.device.create_image2d(width, height, format)?; Ok(Image(Arc::new(ImageInner { image, session: Arc::downgrade(&self.0), diff --git a/piet-gpu-hal/src/lib.rs b/piet-gpu-hal/src/lib.rs index 3ee72b2..fab7d65 100644 --- a/piet-gpu-hal/src/lib.rs +++ b/piet-gpu-hal/src/lib.rs @@ -91,6 +91,15 @@ pub enum SamplerParams { Linear, } +/// Image format. +#[derive(Copy, Clone, Debug)] +pub enum ImageFormat { + // 8 bit grayscale / alpha + A8, + // 8 bit per pixel RGBA + Rgba8, +} + bitflags! { /// The intended usage for a buffer, specified on creation. pub struct BufferUsage: u32 { diff --git a/piet-gpu-hal/src/metal.rs b/piet-gpu-hal/src/metal.rs index 00eef49..e3157d4 100644 --- a/piet-gpu-hal/src/metal.rs +++ b/piet-gpu-hal/src/metal.rs @@ -30,7 +30,7 @@ use metal::{CGFloat, MTLFeatureSet}; use raw_window_handle::{HasRawWindowHandle, RawWindowHandle}; -use crate::{BufferUsage, Error, GpuInfo, MapMode, WorkgroupLimits}; +use crate::{BufferUsage, Error, GpuInfo, ImageFormat, MapMode, WorkgroupLimits}; use util::*; @@ -279,14 +279,23 @@ impl crate::backend::Device for MtlDevice { Ok(()) } - unsafe fn create_image2d(&self, width: u32, height: u32) -> Result { + unsafe fn create_image2d( + &self, + width: u32, + height: u32, + format: ImageFormat, + ) -> Result { let desc = metal::TextureDescriptor::new(); desc.set_width(width as u64); desc.set_height(height as u64); // These are defaults so don't need to be explicitly set. //desc.set_depth(1); //desc.set_mipmap_level_count(1); - //desc.set_pixel_format(metal::MTLPixelFormat::RGBA8Unorm); + let mtl_format = match format { + ImageFormat::A8 => metal::MTLPixelFormat::R8Unorm, + ImageFormat::Rgba8 => metal::MTLPixelFormat::RGBA8Unorm, + }; + desc.set_pixel_format(mtl_format); desc.set_usage(metal::MTLTextureUsage::ShaderRead | metal::MTLTextureUsage::ShaderWrite); let texture = self.device.new_texture(&desc); Ok(Image { diff --git a/piet-gpu-hal/src/mux.rs b/piet-gpu-hal/src/mux.rs index e4d7937..4a54e96 100644 --- a/piet-gpu-hal/src/mux.rs +++ b/piet-gpu-hal/src/mux.rs @@ -35,6 +35,7 @@ use crate::backend::DescriptorSetBuilder as DescriptorSetBuilderTrait; use crate::backend::Device as DeviceTrait; use crate::BackendType; use crate::BindType; +use crate::ImageFormat; use crate::MapMode; use crate::{BufferUsage, Error, GpuInfo, ImageLayout, InstanceFlags}; @@ -264,11 +265,16 @@ impl Device { } } - pub unsafe fn create_image2d(&self, width: u32, height: u32) -> Result { + pub unsafe fn create_image2d( + &self, + width: u32, + height: u32, + format: ImageFormat, + ) -> Result { mux_match! { self; - Device::Vk(d) => d.create_image2d(width, height).map(Image::Vk), - Device::Dx12(d) => d.create_image2d(width, height).map(Image::Dx12), - Device::Mtl(d) => d.create_image2d(width, height).map(Image::Mtl), + Device::Vk(d) => d.create_image2d(width, height, format).map(Image::Vk), + Device::Dx12(d) => d.create_image2d(width, height, format).map(Image::Dx12), + Device::Mtl(d) => d.create_image2d(width, height, format).map(Image::Mtl), } } diff --git a/piet-gpu/shader/build.ninja b/piet-gpu/shader/build.ninja index 6a59f59..e79908a 100644 --- a/piet-gpu/shader/build.ninja +++ b/piet-gpu/shader/build.ninja @@ -58,6 +58,12 @@ build gen/kernel4.hlsl: hlsl gen/kernel4.spv build gen/kernel4.dxil: dxil gen/kernel4.hlsl build gen/kernel4.msl: msl gen/kernel4.spv +build gen/kernel4_gray.spv: glsl kernel4.comp | ptcl.h setup.h + flags = -DGRAY +build gen/kernel4_gray.hlsl: hlsl gen/kernel4_gray.spv +build gen/kernel4_gray.dxil: dxil gen/kernel4_gray.hlsl +build gen/kernel4_gray.msl: msl gen/kernel4_gray.spv + # New element pipeline follows build gen/transform_reduce.spv: glsl transform_reduce.comp | scene.h setup.h mem.h diff --git a/piet-gpu/shader/gen/backdrop.dxil b/piet-gpu/shader/gen/backdrop.dxil index 4ebcb1cc55f84fabfcb69250b09215f2e5f48a4a..10e1bd0ece9195f54e1875111def69587d416c35 100644 GIT binary patch delta 1082 zcmZ9KZA@Eb6vv-?d%3+Wx8>r-J_3)oaBa$6C%v4q=RhP-N7Muy0T3$>vOFVW!4 z7C+E;kaSz!-N+&{=`w@Ffhi%IP7>2noH$6fF*4cGU`1>?Gm|+@wisi))oG&V%X7~E z{Qe)#KTo`DxlC8zY3IQI3KQqXIh`QgBi#6D<@1j|cy#MB(JQ%F)l%Qt)EbZjz=J$+ z9tR*d0}+R!V?fmHTzu0&18{c!2kB!B-$@JK%7O;vmkw>v;USh@0@>pLR>8Z^gVTM$ zv?{ueApl+xu1L5MBgM|!5G{8OtBLp>MlZA2md&`DNZ*OTe2(TgP=xpSVjb#0-`yC5 z9OTsdRWU=P23{wCI?$`Xsh!Bw*wWK79JE1?I~h?VN#AG8AlLpXxb$Y@@P+Sxxt*CH zA2l4xG)S!>t??^&`K;o^Ed6@c<#vta1)5DJH}{|o?mnDMdu?a{?X^&TMG}Q==wOd? zzJIu>v|{+eu2ZB!}Otn0)7=P*gEFxuf-n=)IQ8Za<7}uBEb>DWWKs{xLNFMZOasX|{MeWnmE6 z^&=Gkyype<*9&Syq?D8a9}ngd^-=8UKw`{>d-RAMWL1ZwWG~ap13g8CirN1Ew2quQ zR?L=?lxSuW7=njT*hsLdQ(>eBlrK>&b$Vi$PSoOWa6(k6tv#2pOWn--5F`~XEICi( zs=`yG!;5RiyeIzltV;*=c|$owGwtn4lX^5nt+MnuCH$<^cAr;T_r&}+17LbU@nHaz zAC=7N3-~2~v5=nX*pM*PW4_t|#>XVc!U9P6zt;TD)4n>?MtS!B1{!Dxhy*_#bTEFl zutAMSm324lc5DUY%{m2V!GI*8=@Qlo7$z3%I_qjItnKWHj=q)k_ee@0$xpBr*G`^E z!C!H-YBe`JiKBVnHZE9UNv*`Z#5pkYOx`x|X+k5>Q4h|9iXuEcRU|K+W{Yp*M?gZ( zHDx&#LrKNU40%ftR2@t!`!V4p)D*|0CT|%ZG#Fb;q9FVoS&97jLW@*6BQx~~aj3!n zz=UsLyUwbPEvB_jl>=NgW7Q6fI{ESLpwzbBBiK|r_0g?zkpJZTrstNn%6xOeph~wt zH0O~oZd!W(sof Z!yKHf%faZ~9BfjFu+8YR0|HUkscB-qi0DJK4bi+@aGjCaA9`26!c&o`g% zljr+;UhZTDGgC{gMw9#1v74_Gk20)>R(`VWYmbTm0CXUglkPw2d#yZ9<=s5-%=P0+ zmVLU^T3OX42Y>+*@EHPtFb6zN@QwyvH=2SGEe^oILxx^K=}wHD<2e92NdU~;ENZgrAU$ayarwD6H5S(BU`-k01jXt^V~wYdfsARX ziYR)+xm}@7iPgY*Mgx%|@iI1d4KPj%p1C&na9s!23k~493;`|Vc->=~SoOAK`{jd$ z@lCe&uE^T@!;Q$+;`$c7ur#+xnm?Zk1-uUYHa<4WT4Z*Kqy>H>Lw2>0w-WFo{0zy4 z+y*=h@Gr*$KTOSzT=UK@n`KUY!#jB$jThURT5ZRXVOi?+t7plN&wjZ6e$a1n|6Fi6 zrTmoc(eLG_7K*ay#s28*9q)tN-`{=D`@k$4DdD}1w)TsimR5d8ZO0LX5Wku)!^?2e zElEMRO+}e~aVQz}c!Lo9G1NyxfLUBSy%hx+m<%`>5D=C5L}{Lo$AGXJl5Vjmn|#{o_VUJJfy;{5R` zc3;RSZK6IRz`WuJjZNZ+GIfSyIi=maJVL^OflIJtpR@K2^t3wrud zGM@%{gxGag^*4hf7SasalD|MXP5^O;>5)`ZE+)A=3GwwETr!r5?m@|zi)&7}fJRkk zxU~cdosraX3PPRu93k%)&k$fJjmo%d#E2neM30xYxL~38FikUFcQm+%;|o890&1zV*87HRjQ6rM`-U$8xmd! z2*+zy-A}ZV1e(HM#aeeU3^VtFQU`e<^_Z4$<)S&_`tio`ZW^D)za~7P0c|>x|G3~2 mlLo7GcL*?*Ai&QAc>bUOlUD>dTgk%)9g;7+^qt0k8vg)Ncy-|b diff --git a/piet-gpu/shader/gen/backdrop_lg.dxil b/piet-gpu/shader/gen/backdrop_lg.dxil index e6b2f1aa3f98102c410ab7d255fec28fe675633d..58e21a21a25f338d85723057bb88f2e442e8817b 100644 GIT binary patch delta 1089 zcmZ9KZ%i9?7{`CTUav>FmMb0bdRbFx&Ncyx}M})_D_DvdGUTY*zvBjuA-_>1pp5e zU<3g`nn4ODQ;znO?%<-U6am1;k%!7BGGsR){*>mYG`C&ZewPQGc?l#g0J4-^ejZ$G z2FAsd>pTSDW$~t*>oVX>goT8vu~UaeAJFDZ-fBto8 zYY_Xb{MD^;rL8C1@SWQV$ONoTBJ=8T<^8F(hGlWlc`wfAB9*pg$;wRtA)isBb&#FP=bL zup0Lq_!E>9BtT>Q2Z|cn$L}iFA)$eTKRBGS1za)V6Pq(yo-IE{WiP3yj;TT`AB%B2?m9OFv z$H-t@eU!$2RQt8ZlA6v^nMuo!eU<4PBm`QsJLco^@=_GQ$s2BSPle z8zwwmiXxRnOUdj;5hAv0^Jwh4M>W4;I{fcf>C(nt{8tN-%s+f{oe~%o+r%bcO%a-{#h2JOBUy delta 1134 zcmY+De@q)y0LNdi?e(?0((Cx)byREzRMIeZAR|bo-nG!NNbz8Xz+~n?FoQuMVC?|t+6 zzP$HI-crU~hM~CHW`A&IsN7!XBWe03)A!56Z)5;~;t4}Q^oG0I)y_h7pkx_dvvPv2k3?ZT9e*83W+r)C1KYGUU>(z>jK=8EdS7%2KLZ z!L32?D>Cwuoa#cksK^1(MF8Lo%*gFM7D!Fe1g3ebBNd%9Szt{GD}ch(-L{(kIu^3_ zVbyMc|5ag4f0sceN`sgY+|iR(J70IMaq zh?mCZ0Cpubac_EN^z7iwic>9E8{Wz3Y&_Y~)aE*Ze4;j7e)kyh;j#B)=fgp}_hDXN zilfT%!|#r&`GQPp(Z4aXFnIgM_g}v^c-yHSEfTwpu2Uzwm^Sf79Y+wY^!y4@!wYcA zt4u)xo7$3|#gQbr1{bQ;KHz`uXV@#+K0y34a|c8;9aRH?1=DoND80uQ>8BCSL{KhP z-yA|4$a)r-AUba>ZTa7lkAyTYrj26cP5Er-F2INsdHh zIDuKVIb^qR<@l$8HeHKcmhxs`O-CYOPXNb_yPqQBi?SU%ZUzoM76LEhdwbm33wDOj z9#j4Vgk;no?Bfv9o9bt=&r>z&LC$y==>!pt`|x8sB_n3Fg9bXbr2fcq8kfE#=$!zMY&>XN{z>9@+Cpw{sHCv)4yRxuCHH;&`Q@ zGBFf=lG31OzJ7EsaVx%S$3s)jLOes5lswcRO~XiGhGkC> zq@C_Zt=yr?i$U)sr&)?yng96xtMN?Xs>Ev&h}`3n%PafaZ5DV9t|K|?fv7&Gc8LmC#37&mLf{D`- QoTV3G+JfXtAN=3(Hzz=H&j0`b diff --git a/piet-gpu/shader/gen/bbox_clear.dxil b/piet-gpu/shader/gen/bbox_clear.dxil index 9ce0addb78ff37cc0ac08c5bf861f956dbca95c2..ec661f81be0e0643f9b6ed85e6bb99ef2a302e78 100644 GIT binary patch delta 444 zcmca1@jyb-CBn&>0Szz$`3V~(evAkcXB>(Pmza#A$#*iHh!kbD$EkHd`APsS_2w5rnoVj5aYNi$uP?p zY+a=U1B3nKj~u?C!Yzz94HP&H9cRcf3AZTToFKsM%;W6LAkCEC6ksXb(sFZyfPk}r zv#_h+2E_?x0_Ry6gjpCBE{aWiX6CqI7U!V_J_2p2hbQnJW^3AEkn0>6$*?taawvt6P*>B;q)Fh%ij@kH&})vG5bzM^kI=9LadrUnHGjr+$m{4ruNCu0DciYiNh;}TI9gQN=}z5G%+J_izIj$5k2 zMEs=+d@rQNUbi)ciS$dA_&r$I_*||sKv1wp!USmG!Od*@F&1YjRRg8~p431N)((eEmWG_hfVqqdESVV?DkT^g>?TJtJJ+jm9Ny?C z!NZkYEX{UAL5*Xwb3&7Wg0zrRWJh8GM}qSOAAwbAO%rsa#U4d+Bqbb#Ecy!CUG8GkR!1n@rHxIwW+xpE4XHC(aOmw%hk(FNzE$>Zd$-4 zIj=i*=AD>Zvu{sn$Z%?Vo+N%OE!;xQT-D&!goiuNENq=P^`zz`Xf7%S6V>E!_6?C4Cc8uXH#vH7H1E z+*8hoHcd)kP~c!=kubU4!3YvHkykOfoZ-a4W6r}UdB%vv$snl#q@q{C=V)@T8jFME zged|JtP{fo4t6*)12x#BZ!lzRFO+#c1ul|tz|?ViY3%naaFNUlmW}0X^PbB&y1E!7 ZT>u%_FID0HVd25=wHI|78-TF^0ss#tsty1E diff --git a/piet-gpu/shader/gen/binning.dxil b/piet-gpu/shader/gen/binning.dxil index 50034cce738c087ced84608eeb4c0ac98cd99046..849a59fc82b517ed778fa90ada4ee59404d5de8d 100644 GIT binary patch delta 1688 zcmZ9NeM}Q)9LJyQ9pzfCa?n-}D|iJ##4@kQn>x0ZmrikL$8c6bQ(i@^*g2iANv}tN zmKVo%sHjt~lbPz6!O7GuEa};fF~x}qA?VNm8g;L6ZW2v4vpXd-+5L5SzMtRoe4hI~ zPtKR@OIGDtvU9-yp0g@9ieI=`xaDt^_e@6WW7E>@+T-WC(AARAKb_0IgV}woJnZZE zSxsJ@+zVt)^bZlVM7tTG`=db02H@_1$Wkc)y9fZ@VSfrMa3~S$O^TH1jbvS`AAz{^ zIl)VKU=a5=_z#gCP=;}^M-c|;5q;`gDo7*Q2>@^YMl zcVrkdVOwC#X(VjKVm(@?2oNnsq7C1(r~F5%-ATCk1|AzAu>{orrI8%uw9A>-iWz4UnUaNGD-BlTx*?Ek(%*0Q)_=PU13msV76E!a{~X?vUL z#G;Nhm)saFK00c-G1_iDozrl?@E&HGdoc9zfqSNmjQEjruPR&;#RUieC9+nB3!h%i zUJW-npp9vULRkwiFEjy$bngIhQW;>Z-Q&*mK`6S5cN--%N!WIEE4EuhZB&^Rkf_v? z`0D#FUw2Stv^aV@rjw^bYpBCGNP7V3%FOjyXoJE{hfBGpC3GIvH5kfwLlIASFr=4u zyLt69k}#!{5;H;P_k-zre|91gkmizd*t|T-A1hr$)LTbrbZIF#t%of^L)fP{roo-_ zFhkY`&`j9IRZJ-yYp92OC_9QLIhg1n^OqPE}0H3yt$FA%xxl za0CvYRiafetBC@Xj$HwnFdR4#-mhu2SnCYaO-czjZ3!I#Bfv;K9B@v*P5SDuyFg;4 z${GZa0mpi`T?|=-+DlWlV8zDCwyJB@lkTe4e)wq%K>(C-mWQ9lOF&!HUh;SZ3=6Z^ zA9>YG7)ZGres|%ECQN(1H`@*wPP&=BCJY0Ju>8|d7L81D+#FX8^Ui|60)tWQ4sNQGdRZKkEj!;rY*5C#7BfaQ{0y?enLx`uxUa8*y`D3Wu}<*WigV9 z;203x$qrv*X#5n@EqG~hcR7Wl8;%dFo-S+X>?3G9@ zG>82ya?z?TSH5`>C=dEYIHXkA$oh zPYH#AQ?MW^Y@8=OVMo!VMDAxTyGI#?J`v&k*){CPN&*f0LQI}%*GElZFYG!2C_0Whl3%vv2*cc~qHQJmVvq{&E#O z{I)t96pbm({Cc)jrB?bvR`yCaxf`m3eLCYpv)3bsS&u5&k$-3I-Ak!rABrY)XIvnO zQhW~5tU*pZ0e(6oJD`&ccU$VxEOE*CxPs7EiyfdID6VHVb7^hE)vd$k<`?a^Tn?n> z*NwnLKml!`c^#|@iIbv&?wrLALG6BAjt z?_^onA@u?WQI{NuY~c^^vOen-Jfyt4l4JBEDJ48%i|Vmr=dpE^v{*`^140@p4(7)S zC3VXXfwXu;Ko=pDYZJt({#*fWID!gO(Ml2C{TyB?!bGzOS3HN4Mc9`ozzGc(L>iyz YjIzxuiKAeN;ygAcLf$M!|KCsOUqhujTmS$7 delta 1688 zcmYL~eM}Q)9LJw~;kcHoT;K32f>#iRSYkm?*u=HGO;NNls1;EUP-OD5PMn)!(rckm zEGXOQ#>BZ(oOTmURfaFo+&KAT#hwkyts}0#Ub>TizPD@t=^lI~ z9R9$u@F>D9;p7qL6R3ehq`ZU2$EHF_`)?J(pRQG;G!|iNi(@*Ej?8R5^>yoD8+ES)@j2+ zD8wh?rfb4TW$ zO6!{eEo&xjSOy4WuAfp1Mapy;I1*8p~N}_I|9_&xThU_mLoD3^J5kT-kWD!p%lvynG&;PxK6h`+bn| zxgXp%V?MXPQ2Uam=FqT(VhaaARH~4Zga)uYT8>yeIw1~QMY)&|i!Uex-((@ypef=j zl87c9GYC_574T3sBN}HJ5@&Kj-X!iGnQXrl6EZJvuxX1Yd}0m%Kr1l*l~1Q*4=LDk z9UuD*0TY_whKp(Wn-kh8oiK$Sl`yT(Q9{aTstzA#iAtH|6|j{m!r3?rR6vlxj{1e@@KsDB6IGyw`CV_ZN(2RF%o$DpLJB;sCwqEZ0|>2v{4kmosNpgV`M4l(CwhV zR3eDDitUMC>x=+t4VqJ_g%$@Y6An$Sp$Bmu8s z<p|x{urMP zn$BHQ>7hLC4^`H3U$D@U1uBAqM5}r`TKPL>1Eb6>ZOb>-n9u+i`akXSwen6_9(x6j zN)3gTNTnSbzGT7tL_O~>dvojyHXXrVldCP9Bh@xbFG>b`{D<0{4$y@^`Usw zaH0t$XDW`<8BCCuPJqD9%?lWucYGz?871*4g_MFyt&%!m9BaRs=@Bx>yEesLi|2Lw ztpg9z7dM6wDxiSgWBmq16Bdu*qT{gRYV~1Pop41|8dT*WS{CSU*m6NEW8pHiw0%^c zfco4c^M~SjKxxOEljfqoL)d9zvQ{%IglUvxZU-b@@aca)j`rJfGcSZwgK5^~%qSmC zd-?3c`^0GMQB)p4rart(l6uH_zxf*Jt_bI diff --git a/piet-gpu/shader/gen/coarse.dxil b/piet-gpu/shader/gen/coarse.dxil index 16d47ceb1c63abdaee8fa52a8b5a604855d036d4..a3e34c66bc3f07a896b41c4b1bb84eb997911c30 100644 GIT binary patch delta 7096 zcmZ8m3s@7^ww}BuAtVqWf$)-G4DwP3M4^B%A%YMQA_fGcGJ;r&uc)-OYRx18M2ms~ zie5bi5ex_l1u9q{2|xIxbF*kU;Yu1Qvc`NK>q*l5Q-V)!D%@c60Ov`lEtE5-^I)H}SL9raG{Kc!1kNrel!@6ie)-L1V8;vDwBdl&iHcFPKK&@M% zw^jZ6$HS+G;Xk8RK8>RL7ub9JGuyj?eV~E0b|fo1iBz{SKi|=Kt8v%DOR@Hhk;p&!fZFZ#}J{u5k4aJ>nl2+#agC5D<8=-OB(I z&Mya>FrR_DG}p6DWOgYGL3Y#$jwKa>gr!G0Zb}<&qN45O0!H#1}vewFhB8yJholRLrY9Tq*1vOD&Ky!G z3vDnc-SkedvTzU9k28kXKy_@C^bo;2L#al1F;FPmY_e=D5YHNQ=PwkHck=Oln&tQc z)~qeohDMT&$q4njl#uJ!AbWQ)3fP6H8>}7GaxpeV`a*(iK&g`m)J*L_X$ni8E0w}S zH_LMf&s$}+V&GM8L^Z55C?p6D{&T0Vm8rDNETbJOa9kB$B6e3oMz+eeM9fy*_2)z) z88{?hz{jNh60dP}7}!gc_fGMKPLku+D;I;!)+Cv%r$Mpf8XxLRKX9s%lN9qIB-y#! zSc+qaLtw&9+Py0YkJc!mbr$5aJZe)1Pgi6GmE_Pgt4y7~V7Cl} zCpXaiYw(^5z7U2%SZRwQql-FyCErn7?TqZe{sYWb%lmoWXP4Qq4G{&OKr3lBpWbit&I#S8iIZKnr8(a3~#Narz!xX!RCr6 zl!O#GU&2^5&P<_BAE2dP!g;L61}73t&)+1_m<8IoTI!ELFyih!C6EDul9-Zg(9ih( zqw+w;brV@o$H3a$``JiER>>(XwHM6!<32H`LR13$cHc_Zy$}@gSijrMpwsj^(t>l; z0h~$IL4moQ5DDe50^T&Tn;E&80q#%|GQMO-$u$tI3<|Oo+0W+_8R_rV{8(y;h3C19W`EXt0oWxXr0 zCF`K@7WH4a4TTDck?y47usC6V@%=o=lN^qj{dc%;-N%?D{2$d@oAu0zny6Z!<~S~6 z;K2hD%OPvuUl+9~tv;Z3h2Y$yOPy5v0ZnjVJ?jBOjoJBielkIW)c*9BSH4GxD)-)K(vyj#V2 zHycTnR0a0gzskRf)cJ-AB515qj4Xv<3QxbGLiN|K zUN?{m9JOD?row~a7T$#iTeZv3MzzvCX|V{HGqR(JB+a*x>pj*Rj^iN28nn`>ygF`9 zA5+{zbtuWeELgpz94BL}K$(Fb8_124sdT8tcVr*__536<6Zad`(lU`pGQODBLKamp z!O=Q=YC)tlg-mYTvL4&+pioXpG>#u7K>641N2ZkP#0cyFEsq2>ltGYPljXVN zH_wjU0+CyfK}WxE^n{{l=VnHZGvxx5_V|<;cHgp1AA`;X(T@x^c3JYf{0qf=>NUce z;OGv*n6HWtNFRn)e>_XjE{%ekW-Q=wyc1QscAZUaS?4`F&38-qjM zSS@&Bh2d6L(c75`~Q zN!Y;+GjK6sr>H#iwS3ynmCK|nmnm1;DOW0ubIU1L-jJ`{Azx`FUpYa(l25+!8|g|P zuw#tB`CoSa?aqsY|FrX8X8!FCeq9M)bLRdoYm=y~B0u3?mUTtkXI}H-R?1$8m(yvX z+*h(GAI6=XV>ct0jl))U!Ld(}Hyr_%D|}+lAo(3F^cAkLW5`V-*SeXS0~(y|8&H@S zlZjDA{&DbQK^N4OE{vd&>#^<=Js=~Z9Jf(;F2NZ}oZw2Q=cYS1FR_1_YH&6PMO2yi z7nG*0GUwOL1HZJz3WAC1qWUS@;v0>T{TUJh;$#dSb>$LaDt`$Zm z{HoHmkc`!1GPj=&5;-dp@{Y6QieMr|+72Xj2T!Mna=az(9<^yIN7}5+IT97tNLfhw zjabd6S1G)HFd=Yy!bES}kJR>$IW>*H9@eH=Ncx@VvTUAM1(v8iL79}ec8ZaVM#8CrXRv=Xk!u^K91r!uZKtSjd z>)*TXOzQ8wJ0Aacmu;5;&ClArRl9Qwa$82v<&^nUSj$Vof-|}aE63lTzW0dw+!*q_ zIqW$lm}dEu!mBSnzz+bmpY+f0=cR$-zXE=H5cZoexbL-9mu;^BrR$aVRVjK^3edbK za%(QSRe|r`hIEtoxtP5Iv)_i?#^@UbT<6JRz!V=SaH8H?j)bS~`JneTx{i?i22of@ z1!d%GoNg!-Q-EU%4|#ai{K0sONB907w0>TOg{O-D)|+LXx#87z!}lDA>$UE(eKD=kyDa?7wioe#zqt4Qh4l*{P<_!LJZYnW zL7_G!MJr|WA-6w=+};nJ?61vKwvhi;nPsVkO8M+Dw}oSFA6}mPv7OZK^qL;Uaw3ZH zPWKOMPC$87C5PfQt2Rel_kN4$?A_6Y zdC{I<vK_hd^eiy@A=iZ=KPB%FL zr`%szCQhe^~$6IPMMpvf3gk6wm%w$Rlt zn`o!*77t`A)q&)yz#UbA%9)j?EEcGlKSwY5crcXez|oU}1IL2)RR#H1=5&Bs60t^0 zF`~Nx^jOGaPDZmQZxl~v7wo$!tPlm|n>D7@7AK92SA`N$EOa_13ZglG%Rh9_ z9R25qm-y!4hg?*)o~i_P-^(SbDmUeHkT+4KbGd~HE`^6Y_x7+KK^xe+0ST~)HvP^a zPlF7VlD!VvdOeZFD5=fVGayHUPUGcjjL9pqm19yf2Mo?RYRjhS0+#CXq`F?of=*jG zE=@IfO zh<)0{KCKBpdzag6nO@Z!Tc2k#lXQbgdTv;nDc|d2@D)4PW7zD~9TrTyNiF`zu9(>b z_s*HcyTp-?NeS-lKa?a}I!&d)eF=_AT6+M|pY#KJQ`ZL9eIZ^8N?m^hi(oj{eufdxoPxf`xk$1th$n9=sJx9K&(@u`Yn~x(c=6Iq z3lf^RWn8B+c}#iPTplF>d`nuIgR+ksbF*hnf*>owBP-#D03dicbZD2EU7KqU9TN?o ze)a5h-011;wN|A)OexbWGuD>YRST+r7E~vASMO3jgR%9VUi^>L5O*eF^su_zHYuk% ziFsV=?^PkH$X8kERWx6FYKji6)~XNpm~(AcUG^lQwuQI0_N*Rfekt|o2CJOZgb@V~?}CjjTy2NNvp5#;rXE47e0HA803Qh6=d%z5I1(cS^x3eZDrSQ(4n8=G~Th z5c|yHR^>oujh6gdltjqbHx{fa2f{`RKIgT9_BqLj@>DlLCL*56a6F%j^2Ko;~C}=l-p;mx3FV zDe6!8y2jjqZkm$4bQhwbj3h`~pPoi|s%poarz%a{|l8Yb$_Pe^N|v13f{5ua%; zZ=J}gR$t9|KYITCzM8tPeJVsF;K1--X^9=Q!~yWD{6Igpn&Lp=CJtJ_p#sBPhf`c{3N{R4H1dW(6| zrVDb*GpvXHa~>))<~(Gc1rp@$nc?{#xzJ^ELsFb;!<&^2eW?wiq^!onNsXyyo;<#} z^W2=l+t_&AcS>#+ymXjeSWQYZKX5K<*C9tTxNG`=_Qp;ib`p8;)R~oGS#kU<_aqGNZ0bZS4z@fN2Y^K``X~E3IRKu7Sf#<+|2{OCi^?LXu9BM z*i{lCh%yVh{4YtI84G$H>0lte&|Mg1wpYP0fE@KtDU;eK`;#B*9&kAf=@%er+ zO?}tDSHdxu$itbUs8t?^Gp6nv1gRuY+WL))rt9HL{=P>K_crx|+Q$k!o&+v=sDC1| zH5PC2>*7US{w7HnmDb9aLzz<>p=7{m6^4q0JO_!H6e&;J3ddnQx> delta 7027 zcmZ`;3s@7^ww_5cLl}|}LW1F;Bq9g_bb=s?4>F*LfH-0VM5_Y`MMO;%t*<&sc!(`3 z%0t>yYY^0^peUf&>M;m6TuJ!^__Kn&7Ho{kK_5B_ri9i)AUVw z=?&ZGB+rk#B!?ghLZJ0r2x6Wg02hKUYQX!5NYgKJ9t34JEU*t}Aw_WIjMUQo4rdN< zP*3eOCoN09hNCsPLZT0$tV-BGq2N0R+A4*h_|gpC0@@vImHBWft5c;6>M6`8UDBjd zR*x{)nWG%^;U->4m48(rLs7kqiWIJudDuff++7hz0t0TX)#GR{2X6D>YR{*d3RBrc za?<%KgM#XuR&(+DXIs~7jLq9#?Z2^_?HcJ&njTz!7rmZ{2E7PsJjaUg7Kc%1eF6*oH;{62Mz>HC_NK!Aq4toT(v-K9=ZhkI`^-fIODmHuKc@9Qf@9lD=Q}Cr9{3}> zZTJ^LNTMNieiTxu)|!f729|ZkWQX}{Te2b`wFV_&Un-P(L#0*ZY&CQW3&1BJ&5dFnD)0e#~} zcUV-2%uN-`j3(%!8#7XFs?+EToP-T%$tidiPYs=fA7-l@*%D|C8Mx%-~qiIglfpTtoNU)fbvQ zAIPOsWvs1IZ8n?KhVW!cM{_;w)M0_T-QTckQ-!;?aSQoHrg5m@0VSsR?D@Tsc#euQ zh(jTEHP9Mq+;~rYDtmm#F4!36izHjLqSr#6yATfNnF4gcklu(4JF^PuaF@7q*b-Hs zItO?FY9i0l1l}v;QFt%e2D_?deF2t!j)lLo-*FY$<_Ww-*{vcY4Gp|GPBvd26aAaS z{#ox>Z4Z+LvzNamAjw0E#0faDO7ezn1Nk_=w_@gK;Bb|b4E92~APNk3@}pZLC=vLu zUQYyMJ&*v4RJ2sPRG(&Gk)?z@)4o`T8B+4WqtPYQFVj^}lVY~IOj9}lZ{ffaIEYL= zYbJcI73p9QAin8x(FEnqk^th5R*i*!URYudupHrRa)_vkCC8_U)no_E@cf zv@a-@N||6d;^KE!bgT&B6}sy`hv>Ye=rwMptD7A!8)qk*4P#b4W7NqQ$4yD2AQ>0^ zbXj9-tS}6!?p6g*6;@Xx#JTR0uft3(9w|ap0ku{kRnb)*RI^-@rQjc&G@g6jbBZktOVTWtf*HyS4d%}uIc1Sjx+>betQ1N|q zvNZ2;VGCC=jqfX;Od)w&>Z4RJUg+K?vk8n2zQk0JINRDr1l=+zJ^V=`UcYPY@Z@H` zUfB#4DOl2V+|6uWrfZF_uRMf8EIVE>-(=d?=rC1LFRYlsO+zf(^RE*yo~_#BUp3Xh zf)#E2u}T#rH9@^^X#@&8@F@B4hnH+^U^rz3IPRDyvh-t7sV5R zwP@{~hqy7TYy#!keZ;81)%i6yzEymeYqmKwFu0>uk@L zb|4kqi8s&uPNU`R=>h)_S^k%UK;-&M zRhb{4n0ArwRJ79-mV`^1oBM1<^Mz=VMpSkaL_aSy;BOY%z)KUNgz`Sn7JY05qbyd2 z0=Ej$KN0k*9Zl3y)aN^P0}o(@6+ax@&7#ywHNvW;w4c)#!8UkbSIbm$>yB9XM4OiQ z{KiR*z4;fYHGu6}tgF|vdj3&D+6Z)-)~ptyYeBhG4UfG-6JOdMRI796 zgtYN_p`3I!aOdA9l(1j#^>Pvn>hhtlhhryF6AD(@fvyws#ApVD=0NCr2>lB}|Af#! z2z>;hmm#zPLhB&(ID}Sd1lu)&O%U1-q0b@o8I6-th%%v6KN|W{5c)_!X$6%u0|SDA-vk4Nf^XDV^4kT z|Ep)jp8qvR`BLCNY!1+q+U;!|0E5O ziv)8giyO$sh>x(jzT)j< ztL>a;#y!WcS^{)#pu-vGUNStoFk_i&YUoHsFM49!L+U^t_0k@`7)}SIj=Yc&hHtVf>L%{Mm(n8Wwag zc+H%P7Cwz*piYu@9|?=40Zc9M4nJ?1`bXF7`_8kQr=)2>QKV7P+s62xioP4RDU5_) zWl1ixEqt8}g!=~E4pbDTA17aR&DdBy<5}M9zQy4UoQ*h{`oljw`d4_g2X(_^Ipx^z zY>9E-UTH9+<#_m$jiK=Ce!=T{$Jg5Ux|UZe=dyFmd@G$<%gO>}TV-1Lh;VRJ6L%4IVqVKCMeDixP zA$i7@bnLtPEFAO!cH30E&=NCO2kOW;D}kqVFxCD``yamR;|(s37|aVFw1j#0?n|V? zZ88t+{=&zcp^fMeT^&C3YnXSpWAVt^4@SI74>P|jGrxQK-MdTi$^o#v?V~edkH35D z@g_a|jV1gIdj=3tAN*I_hph&t6=;_H=-EvSUH@e0I6X9R^AoFQ9Ys?%nadN;JNEB( z@!$PUzx#0WI_P}iTb)xL3aT+}UC~D;Pvpww`4PVP*~cfm(=FA{lJ-F!GfmG`9eJ3F ztB}gaDzlg>x2`aXyIC*#;;=|jCHnS|8<6>q%1#%?L~qtR$K*MORXP87XsLc`vKfz> z@&W^GBNFL$hYSGq42zzpz{eV``KMWIg+^XyC;&?w?qq`TZ22$tX<1AsqmkP>nYG1K z-^VOBH_*wPEAZ4~Inn<8o;DeOPj=A`y$2&3znp99a+Y~{gdNuI8t!sl>73XcU|~mZ z)~@b}gH=oib+ur$ahmqLYQVlVE0~r;h3-)gCbUr%a`BK%s!V1K{dR9Icjp7&z26#( zw`7epJQHSqS~aMw2D2Ddg$7n|r=kUPYb*jV{Tr(z=xq?7%P{0y{O9#uPm1>R`0mTJ ze@7Wwo`;9at{WUPT_a$zK6`hF9_x*s$8%%)slUBKf20pXdrjQER_@+tm>#b>)jnJ7OiO%FEymsy@!z!VdlOw4)8-`|35L_q$m3^6NcOej%+r}3@oHCLbz5Z{V9OP9K=2eMf z=qh#OChdX5{&}__Umj3!zyYZi+f@h0ROd#9Ldz_k;Ox!@B%o;sj={xc2gsPC>X_+j zk*_(x><4mG;71*mRYo8+9LbjlY#i`$?D08~gyQ=I#nA-COOk(%BdIx%sErv-rAa3` z`AF>qJG!Ns9LII;>u2YhI<|6Ba=B2Qzyf{ z7iuIuK&6akj)fKR=hacVNqd+h>jj{1_{e&8@uaFn9kC4ehU?-<^W@bCvzB4m)zZmL zi)8VPEMH+B?Y_kE8gW?{bM~l;#Yc`DKhE0U<$?4CgU0RSvKhSwx6=gR)w&XFlvnAu z;`opA7Y??@KMjueLhS%W%HBFpy1!@!vl`?gM{=n&Q5O#m>fAK^rU5%OO67AC`X0`E z)VlER;Dxe~*m{gDafwAw$_^8X_Mk?VKWDD{bB?6D)V3W-2`?wam}ed8n)P$nth$6B z3TG0?ziWVJX@-*zb_>}dHzkDhaT+*v3?Co{)H*KV!&D4zm28r)u59QVDDMTVRW;U zl7!R2#p1X<-O+ox1IuM}LJAQ*Dx4CMP%kbkb1PHNFIP`0hZ08Vdcl5m0OC`C`eY>d z9F4u2FK!pLx7dO*WJ$ug+%MJ9hq~t<>Yj8<)UMyuJmIh{iF3ZK$+Ra{LhIy?T>}a+hO`ZrP{mK*kR(=ilxW!>nBh=94SN`=;?NLrenqLL$_^oRz}8UkX9(eJAT0{F*5am%UaO^s6ngj2oG|*^jY~0(UUs||QXqQ$Y_b?GWTUdNk*LB2C-ju^|Q`5GStnN-M z9^4cRSCgA{-RDkFO?RoL#5GMjPjb3O10XlHX`OOI+)YO6px~vfr1TZdwfI zn881noc(wBrn{z;(RA5jPF#qb^KPrixlyrnNp}h!iwnR-7+a7yjr)KkatNgHsEt+W z{bSY^ETNV)WGp+Mu`GLuWxY9oJ-y80uKudHB5^B{q|5o}@q9xA=Mhe_`WFZGcLkc@ z=_d0uUK#Y1bf|alWj#e7Lf;I(7*abt&^Vye`??r0cpA!e5KdS9N>;V3U`W|paS;J$ z$u+xyj<7EyN&TADv^hy5Dftxe$1}hboIcQ4+W!un(M@j;s)=+5*OAALC_oQu6+9Uw zzVelqLy}(1UD;Q#;!4PH4~us0{NCk*qxYB=mRv@fx2={u1bT*B*#9yKMyXa7wPKrO z#o)^oe_dVKz zYn#WY@s(k=3R-k)OcwI;PIuBLcapM~L;~xqz{AizL3X(y`zrWvMLYV3tbqu#t~W8R z|6Z`UT(J47pn>;S2O7a2KqLy(Ljkf>Cf!qfSXD zb>KSqrEs(iKC0aP{oUOsdVu43p+$)*K=lI}$I*H?vm)(6L0aHcQ2Q+W;}USoV+Pmh z$@*Ewk4VP`-e8F%5jcrPC9RdVA~hyc`P(vc@!Ou_TRp{x?gbFDDHZqw4wL|Xm!YBD zR_dnEPAAQA?UM=f0xe@XY4sT$c`Dv{-Y3E$a|JABDSLv7BV_0@G?xL(92xMC9Dv!; z44B%@fPN%eE0Hujbqq801_ljH zX3$R0v63#Sy^NC{bmKRtpCsCn-ujglS?> z?y0h>T3+jQGSMGU*sfwd#%iSI?w6{dk7@WZzDo6#Y9Q%GV)E@+&HgslgS|7?3Tn6i z{r;&v>ZFLNhH>WcoAF?XD>^f~rhB?DIXky9`2FnM(x+&?I5Ay#>&C6$rQI68aVyrX z?O;8)GT&XeJ@C2Wes^u{o;CQUjdc+@|436YYYe&UHfs<O0hFJE~+o*4uW6a=rVCZ@0v?WyFG_vAwd>Vq)Y?ebi+Y-Cf@~{;e^8_^7X!Nch zq%BWU@7nWP4E-M4wiH?kG=$ZSxEP3>J}f0?!&swlwr9T`r3zAtuiF#qoOqd7B}ENT zdw>}iLB~(Pw5b4N4B<|oWTNpRSD9df&q6$EU}lHKPf+_5uSpS2OFT0t{@6 zJ7H=B-AT)I;OY)Ny$G{Y-UbiZ(znfNw#;}|Fx!GbQ?zLwi%_mWEz9kMz=ii4^GzHI z1=3*ZC754YHnAKq2c(>!e8-ttY8t(%EGYuIFiRCF)esw{d=uKd{1$Rdw2w|4L4IXN z)RMFYzTF`s=(o6Y5xIA(El(rFMy=qyvpJAe@)L z?QPfH$yN(T88G(RS^YrI3logDmwuP+@-FX3OI>EzYv|(ZpqIxJ^w$u%zeGp8+g(}n zCDVf2yhy=J@aLxm@X4-dRDEC+meJ^5Yj1#m@k&{OpVU40vbRPT|B%xqzBVh8y!P(gvI+-o7EaU0%jLIib_}(fX zt_$qu3(q^UpUYIp&z2Wc884owDj2`7?)Z<7z{1NJrLu`0lE9c|C;8i^U-ss#WvP%2 zuU`iJe)d62PVBChJ==tSu;7C-b6-)C5w@6w&iQTlpcPTsO)^n0nKdQn5%?4T9Bc!J z1d+)b`5Cgo3KwY7{3>h=gLS&d66W3|B5#i#z$Lv9v>PqIywu z$!t#1Qk30F(mD36jm)_z(%*E|tNTnxz08}sj~KTa;*ZqBD^UxZ&_@$K=v95wl|oag z7kV(T)ac)w^A1ZwHW<_b9*e@sQZ~pVj0)>vt6eu(A3IR5v`Spm+}D_^Kp_^MzbS0s zSUtMFhSY%ujWzB~^s9jeu@%xkr|K7|`h(^Dp5R(fw+?POLi-2PKixwn+3e&+VI#+? z*5~pN`(2sdcMG38GH{Ub=W_C{noD2wK@su08Y@)lg)WZzpb|%K=YW>+)Qk1cUZKWW zgrUAxcq|Tw6TpU8C@v`f_e_!LY}I=CvGumY)lT6S{WU-#5yD*{N75?fd~9^rWl0lp N>>G%7fbc#0{{{N9tmXg! delta 2058 zcmYLJZBSEJ8b0^tCL|CfALJV}5KI6Q_oB#$OCbq_;ufsB5&|Z8i9$i4gNT%^qud0D zfH($^!EB^8h1xQTPCBkT*tOjRsPa)~K-aFd+6ilQt?TSM?hbWV?d}Z{_x?KPyw82l zd7tNb-{897y4V7fK7XrW^#j+1ccYqrT;43uNWcFqwwB4+c=g_&)Q+(p+DrQU_mTPA zFZ42T6kCJ;QUd*vUa7dptDg39BAw)yXhI!Ej!MAO8j$Qp0DM1cYnzNQAy+a741^uz z@G)@P$(h#U`T2m4$!|N^)A}50E#noIcU9FgiSeLz4h#H2$=q_WD&dn(xx+D-OyqBB zkS9d}C9_jIs8L^Qus_+%$Fv?fxZ$wL9MRZE4R;;r3^{Xl>`Kj_9~>(*J}Asw$voqn zzjVtrRdnt27ne#{UCN!Ut6Qq-TdJFWjSel7d^i(VAuj)|`G%}q^4)QelJEFyVp+yd zMpdRUBXgzi5%u%&OLJ5AbzhvmG;v57^_g|U-n+xO zLL(lr3%5FfoXNNmU(bra9FX>klK|5m=!h85M>P2v3I+q%3F3vVG~)HL`v1#)5o@^w zz=v=?!T@U{*AZ*1WGGwzir)pemjktWuoY*lFQ#Gk^xpJ9kZt97b@*HO*E}{&1M-+Z zbDq(HF=QT~X~;aLoV&D(isNhoo%Cfz$sO=fOnlq$gfm{~v$0bNVgevq4qIJ1*`9K% zeXaux2cw(a3?}~9klL0kNkF63P)CorgZ1j&f?BXn6<)a_dydr^o zHnzM=Kd)~bkG6800rmJGwvM46A+jPYOd~!pz8cPP%(n9PE-n4-PlEUaf$$gB&%miQ^=KiR z1F(SK3a{~`Sw=Hy$iu8@9u5#CSKEsGpUO{>TpC6d?|VF`E=bHeObI!vO^Mow=9c!+ zd&P3ffwqfq1z*iO;xJz>(29mw2>!}hMrh$L7oKxe-YM{hjyGYI{w^eBXw!Mv$L9-; z^q4SE-6k@#NL$RQFNMZ{#tZ5*J^vf#4ll%!OiLwr&DqvzgN}U6Llw<4cCR8XJE7lSB(vP zJvA^}q^S*`cf%J#X+~k{KAyt%0kg%bYH3m&Lbe_C_`7s8N^#PAu$xr1+Ev+2 zLZ9G+@N*onSR@mkIYz}FPf<$6S{_oNMPuEB2SjoN!Vg6W2oC)s8PW*vh@_b}T&3BN zEX9mf0%INIM-fpef+f9r%PLlupk-Y1SEc6rCFkQw4ZU=!QE;dcizL47HbDlK0_DFt)yqTExnVD+c4OB(EzpuD)?oITuBmG4HxxV&4{G-hQOqw$u>U>1978%`QT} z9SuVdTUt(Uv{#0pTVg70k*$U|*;1;(z=S!+uGH7j3n%tUT)^^>ZkeT9wp8~AW$y-$n348!#y_GW zbAnWo#ZI0UHFI5BOCcZiUQt-WQ}|%{i7l+#)zqw-&bc0j{zvql&IKw%(Cb5Cs8rq8 zeWH~0AcT!9_<^446vc;I;jUx_!iFy;B8j^CUtbyR2kI88w=Z;-?edA1=r4vEsR*fu Y25E<=&ut4d=F8U*_zYldCc)D8KPdvc82|tP diff --git a/piet-gpu/shader/gen/draw_reduce.dxil b/piet-gpu/shader/gen/draw_reduce.dxil index 9b1b0fd6893fc7bdecd4093fd10ad0e5e93d2257..c6c962474bc385ae7d5db75c0397b5bab5dcc8d2 100644 GIT binary patch delta 789 zcmaDN_eM_ACBn&>0Szz$`8gXU^Vuz1cYE&tG_%1jEL6N_$GKe&E)kwSXV@7S5*Sz+ zJ_D73_zXZ=0H|XNP~6G+z&%AC28LIg53;{!WHB(XmYBSSvqh2bXaGxUAX{rd1IrTb zmRlSQcO@B488e8a0rgc%FfceyUd!w}c@7s>y)esyn+6J;hK@V5n1opl%v4}F-_{Z# zu}D&e>4t!faLbul3Je!m7A`c?aXBJz>4DE>mK=^1yeuJ%2Ur#gcx-81y3#&Z#lyujow=2@CudhFYbVbZQb|4WA?C@cXZKz`IdT4H&E!H}yZVlU$`z5ONfQ_pIM~ifI2<}2sm#X4(3Fs{ zz>$d!&J;p0mDog8jFK8a>L*JExEA!Pu{f|!<-5$#D8PI70#D=v*3B%3bq*R_VA{d- zfU7~9VS~VqCLr$wlLB{x4#Osa15F0r4m=?Y*ByD7Me7fg2po1`4UpYo5Fv1K0TZhm zj{pP9!GtE}_ne0kOcWWkew55$rtZO1rmym-AHy75Q>{mf9*YL(G^fS=3$1G9ZaGnZGQ cJ4C#S)w-bh0e_$Uk7oOZ=13-XP?`q;0G?L=x&QzG delta 781 zcmX|M5FRDJAb1{ElGZ+~ZZ^^7u4P`ADrp=4fbTyKlUI?@W{j^V z6bz0t0DSu`%+Y>@qMD4v7@MU^D;FfNrr>fcsGYEL4hI;~eNu2|3ZCd@ty8Tn;1n0fdTx{FSl;2rdXLv%dC=qCzXR=~_1Dhs ziN$wcjk{y-(yit=XOA&yJa`#l@M4a9wg5ikxUas@c-iJi6*gO<_Y2SQT-B z;iV@vs|W%lLAYZ3cAjBCC7n7%sF`#B)(WONIDQ&IMV4`TdEIE?Wi9BcH#spJwfeC2 zf-Nn&ln&#M_~Zm7=E7|lvCCBCzmn|FKjfXD&BJU};CVCjC#LEREEfutbYZy<8afhPnHz{&@Nr|kK%(g?AFHTC+;}fbO6e*; zo8t-*NzNM>Ll@%>K!j=4k^!z+TJFxJdwt}~LBzyh+LfeB^+i_DX3>T+nHHCXhH-`GEpCO^ zQXLC5A$KT$kbFeoW?CLZ=5S1iu9k%nQwZa_yh58Wf@oX_y;pC@wrnW0v>5XnmCRVb p?o&vEmU6%4#lI7p981gc1%F&OD{G`uZa%Si9v1CYr{jOW@DH5N{G|W@ diff --git a/piet-gpu/shader/gen/draw_root.dxil b/piet-gpu/shader/gen/draw_root.dxil index a84fd4aa668225d7ed426d81a25bc91b30599fa8..d0a326fddecd7759773f6d6aaa0833a4b3993080 100644 GIT binary patch delta 732 zcmaDM_eM_ACBn&>0Szz$`8gXUjoB?1imYbeD5##d$vR7x-~cuoGv>02+%(s0v2fz!}&hZ>VG%YvB-CKuaUBP5z6 zq>dLh@E+XU$e_S5LqUY8bK3@!?VL<%z5)#CO$;{DY}pcK0xY4i(p;HLN*aUP*A%0xrpU-LW(8#N3*FdrCuwQ`_?-@ndP>YUZj2uO>X)d1hhj z#HlAu_fDL;(U)_QQA&kgQEG91X;E@&aWK;gKFN9BGo|Oo$X$t^C4KkK6b9zaa@@j9 zK_bjv>q?kUnJ-}TmE&El%o!; z3;3PFdL$$mIGPxZxbh#cePus5Ihs$_xrs?t!@xU$Cxn5ekr$}gG2t-8p$h93uvZ%7 zI5-?);E@*aH()rnf{CY1n#qo#Ig%}z!;OcBf#o1WBdgtw$-DVf*jt=+dK$S7On$~E zmtwGhh2_l;W`U3gT;@p^Kt7qQzs5y!7w` zn^#*+Tp~PuSU4CM5*Sz+7?~ItKzs(Easi-@H|z`yPRktR4;(EDg6f7#=b$v}SflWCLoelwe@+ocxp1w_cjXk<8V@DD{D{Au2R-co-L%1dg4ROlT**`y?S!&)U#zQ{<20k(~ZoncG|4i zZeqS`r{+OchRt5ApBN{L@G45nco{Bu;$^;|hB>Ptgp)y;jgz4%;X#gTLF~oJ0lZ?0 z=S4T(?O^s*P>?)mGE0ELw7@#4fkA@@-xOMuS745T%SpTm>K4zOo;l9M31~%*3RsVc;FW6T-lh z$jdC+;23eZqmdElK81A~*sBb392^cZ@JI{v8!$Xu!Nk)h-DJnm9Lbi<;l{(mz;ckG zvDNOu>bWJJ&jxqlV9=4rC2OzVSe+2Ss>&CmwD0!kk=;buW{03+2nAdN#{h9 zX2=9?Yl!$|{S8f;ESnw99MU;+NOQ&`ZUc!tjwD8&G|4F{2O5A0hk+T0quJDcY~%6! J07|wn007!1_iF$E diff --git a/piet-gpu/shader/gen/kernel4.dxil b/piet-gpu/shader/gen/kernel4.dxil index 3b3c42e8c69d2a7f2bbfca14d8254af7ed99cfb1..3e21354162cfcbfc0937e0f4e6e1894c546c9b19 100644 GIT binary patch delta 5994 zcmZ`-3sh5Qwmu=rNdf^A0(ip11Q0}slAt1>Llc5Q5fv3RijO=1sYXjht&d1vyp(8o zi6B0LNC#W%grZih&Lkv41hk=bY<$&7ky6KTkZU`k4}1TEC*FH!t;<@l_Tm4|KKuXn z{=WYt^-CLC423O z+nJE}W@w}t2kk{5NLELuE>ST`E%&)>3Qypk?Zqg4dcT_HBcuw2P>9!NEe*a;c=|~q zhhU@G9O=smSu{J;%%(aB(#mME?xKKw-nWm-Bh1NJ9etfeQvMNn4`J z+sk7Qo!N>Hmp-_x*^>k{vdPfF%cj|1P|}#|8Ks&B+)_cJLd7O??KX+<=>h&+nA7k4 zu4}}7{nFJo%Z-7fR|houma<#CiTiPZ#et%gItel(Iu-DZT}H%W`|a~JAj6^h;4X1<@7zkTOQtr&s%1%4&|@nscM(XvY#tZ9?9jnLF}}KiIWHo?pDu z>w+Ki)Z(U)lk2D4%{?CS_))zCkluM3{vB6WKZ&i9)~7`UEVR`B!M_!_q$S`3{x}@4 z;AucE9S+DDC)NM)hhWiNJ`=89IVBQ`;&&sOJI9O3u$)kM;lGBNdV1+pWbiS9+t6Ba z=P>>JmA&6RMl?qbZT;tAx^2Gmhwl*0!=5wX;`;t|-#tNalm3$O_%Pjm=HWLIu1I3I z^q@o~Cmi5hbI%pIe(oszwF|+Kwd~&x)79*eCsMA+@$VT`|CspaA1v_9;Tsn|pDTj2 z>`GPS#2iM#6i6#kAMIj_OB|%VTv;{5Cc}->jh1I5rICa!pG*suv_ancv1KH-U9j4_ z%{;N}18#5Y3RZkvYmbBflosllGV6*MzFvwtk@|YhTxaRx4a zn62BB4Q|h4Px?n)rhiW6bFdQP(*223EYT^ifE9&tlAEjG)An&8D{M28@Zf?xagh!# zO#L{&p9QRjWHau)7vMQbBcu}W9MT)dq`DFR2-r4n1ea>AqcM0>*wg`DHqUOSg+)~o za`=-{unjF9wN5AeUbl%6sG>(z;n#i>1vj?N2`lTzF-;aGb*FFjhpiUo^+E~!3ej8K zWtNj4_8)!k-gwUU16&*bTZrjk+T;@WpNJk80Mkets)?(z@$^o@O^ybs|J*+t_g*29 zgNua8pe?4y>5I&168}~TF1vK0xf184W+RFejejTYtgTR$1}Tb@h$kHfrTB=#3Y>FT z?LE)Ju)^ssKfLL7c;KGF=yKmmR~Rg5*2TU~UrG(3;pS7#i506x4>lu-Z0#OIz0L3Z zC}TQL!6%Gyp@*UCa2%z~sU{Cgq07lv(!}bUCLG;ih5G$t@PXE3u6PA*aUpU3M13ky z;K>C`)ItekqQERqfm<|;?`cD}!Zhe^I0GAm+}szX$%Og5<20zvOwc22h^)21lT&N_ zEX(xK1pb_O&z^5D5SxVTpdMuZp^gegvfK>6TAJ@+%>hkl@xMQ^5lSp9f1{C8c3gOW9gUe9bpda0f(G zKZsy)?eV=z&Xn0)5T8c%bDFALQ#r;VxmOg6QR{B3w)GGMsnO&yUbyj2exq-gObDds?oZW0o{%)@la`2~c^6*F#9Rwp>KH zGf98RIL#$gNvI4hC~x9k!IVGKM{JcAGbfzOBqncNLf|8T?M2O z0E8xu-$+;hh&7{XQsN?iLHE}Kv^i3R?5zQ<=yPQq9Gn~tFR;T#8R zouJC2nj~372mOQ{s3w3HfvM)4? zRU;xO5>=w_Ay3Y{h{zlC{)wEk-{ow&OWs3B++pbW7TI6baSYjC0VL+&-;^v~mYNl4 zytOJjdY!ddLfo=rF0xCzJjekn&vzVli}|*m6XB^rvq!OsRN-^F(@h4f*=c22rN3g$ z7Unq8?Vbfu)Dpt)$Pt>?4-bdcQOw1hM|EV^+)y4VB~3Yv`e7F>nh<`)%*8Y2RT_w@4cAC4*0`!3OQ8IlSxqR_FZpZ%*AAU zV>Htd4a_BpfaecN^mII`e({U84O+mu3Uz|9XKTEYz*zQ%BTxhhxDtYzQf>UGKc*#1 z{GfmTEfelTKMPcyjL2AYA zYlOWC3jTQ^7dq?LjqtdrAuUkljYVM%Z_T?G<*Jk7&xC1l?A_dM%rB#?@fOTDV!oeg zISqLD7mW&C)$v4BDsW(lph-o+R4i`VLrKqfsw*ob(3bZe5eq0sw)qvMB}|1t|-`u zy51^OgeTjny`TuoC#az3tz{lX1-+W7j;N^xj@_C`*ksi1AE4t2fLf<)cvDUk)HF~E zIfot7c#K(-ugn9rSi$q$bxkczDxXJT;1EUqpfvNgxu5vb52WUSN=45K3WmKts_?C3 zGl3_x)y_&N)_e&{GNI31dB3Uls`)W`XoH%BpSc;yYr)9v;ia!rka9vON1X{CB&`3Fg|_G3>&j>NpkcT;)H@x~nqV>~#2Nbb{X z@A0bF7?NWnSw1;%YEXw>$&oOD)F#GBma=Un7Vb#dF2q8WHsym^&g@xZ9hQP2x%IYv zGd5fIYMP~Q!-jiJo-pX!mS*WSY4s;RmJ2@lI0Qa-JjxEo!p+t}Ht>b`^Nef>)rUIX zD6#jz!qjNERN0VxLjNMPFLY_a?&^HMX~KE85-RG>g_V)zG$}PgFCdiivNqS1vFE5n zsa#|`c@K$Nz^vr7Z+sQj?%H)x`$U5^@-FG4?r8f-(8!oGyCpUV2Ly2W=lfxuls%hg zheOZX!;rW1`j#lCZPPwM>0MDk5>aXBj5ePhl{WL0B8HdCPyQF~&#-(m<+A3G=J$)m z&e+D&pF$<7i}3Fj@8>^k=RPO2a@=mJLs1 zhxshcNW~O)vHBOnWa@0019gEKpQFou99lN|Y}uQkWpBWVWiHgx^wHF=LN6XQgLYw- z(n`1egZiY1fSaeXZqE3Zy~pnfu(rQNvR0OmQa@HtQf5(DZYM2utV5s#?Qy?r4>zU7 z(;jY`Ck*0cwphylU1RVajbw!kzALJc?&FKqcVhXw&{*T+`;D(IXUSaZiaC!}+=YKn zAU&SmI#)bwoL4{~^@Z;IK|d&I(!Dg+gaW8TC8QM;#0{`QA z>DI#AeYPmplY!(TLV~s4iiH_(snCfSvVffWv4f_}Iv4JCW(&cHmje!cZ+S+1-Gf#ShwBDao;WD? zd8?(IS8MRio5~95UFhlTh<}WS#1|{xg`>4G|BJ*44g-6mW&sYr$QC@`BK&LD9l`de|={5+I?3nlr7Kyg~DFH8$Kr2w0j#7Y3 zyCABzm$h5FRx3Bkk|;d(0EEfn`Pq_1D&_l4!YMWNrv7k4<`(gAbw>08PO-hUE_R}D zb{CS4$dk0DHsEFK?eDAH<&os|!Kw3#7NzbUs#U&lM2&k6!j+K!yv zqVm`pcB-xYw4)abrE#D8W8wRsejb5^AI5bi=q1y}k%6USysnc&a@)LjHC)wBmtj77 z?T{yIxK`^4gT6dIeXvq`1Q|yEd<xO_UfRZloO0i zL~!Qt(KBb~^H2X|G<(kW#?_mM*x7zID3|!hZ0vPoi)9sU#ZRPtQGwn#&aR?qUwea- zQ;n#>Q5CRo3smk)&Ssy(fj9Cu>a(!KsUF4G216l4)i+v8krHFY?84cG{m7rLh`ZaO&IYYk zuUTi?C?)P|%O4FqCA}Lc3X|TYi$Vqbsf8%}o9GQD0R? zfOL+s+;&u%RTDRk%vv+nk$+)GZdq^XIdQGlK-@)SRlUqi4SE!)#2Rog{!uIz*8Js> zr!7tKkAoG^jBzJz9Wc#{L$k0KvI9TeQiI!F%hans63_I_?DL9a)%q zl=NHFB@!TA9mnvA;g*MMlvMhV+|FhR$>YdY;;RmRoq=*(sc;O*ck}H%m4sB$RG>MH zflg7XK$i??m%2jGu8fqTExICs6oCpVZ1vStX)GKavtE!LPn2^P7X@S;MsnCP>!nuO#On(`#9oZc5Z4kr?~nvR!AWi_U$%W5G!4{A2_DX zwGpNN`nRJ^uT&eyc=*2#$wf-f6bB`J*^!660wh^IVNh~khYkJ(DCwstR>+IR9(!~$ z4IX=ZwDw)-YBKmH4jltqjvsfLHjQ;i|K^h8RhFLeH`=duTHkVx9zd5={Vg|se$i@1 z0^z;81D#Fx5Owfs)8;kSy#mTG7<51IKFe)6ex$6C=D2@mYQ~ZK%<5qD>`BD~+nmkU zL3kPXGHg5e&4p{o_auFMe@n1}nH3?{&A?|8wwm&@4@!J5Q#thWmIh0dc8)W_oNj`1 z;eiRTDi}qtiv(z$0!23-M$uPLbog8p?G^!a`I8Cs6cpS$0p5p##Re2SIsr~a!Ojd6 z{3{AJvrzDgsUWlU5DG4=LBa9CQ9-BaK zLeZm((bOd<7)n6FY0)V777C7OLcuW%G<7}-Hf5pUh6(U26fBzn-$udh6W}Bitl6wW z(QXtCn@}`65k=ckwC+BN<~O6Mx1eb0b`rLfj6>y9B^~)Xt#uxQVYb(LetU8NIyV delta 5973 zcmZ`-3se(Vx1J=EA%Or1FDC>{cqs~M5&;1^?|`ZQhrM|vT3s4FK^zxTlXa4NVX8vH*l-Z&IMfuIVf+ReD5mz3mue60v}W{mpc z^gCfaA4e!Vj?)()${TcW$0sNRNo#48)d~h;X3ZgG3MFcGG{YlqeVJ53lEe?DQahky ztc~}zC3}3-C`wW*$$0uk&*I@DEXp7(EsG*0T#?e7OBxYpHS0q)JBi_wyn&RjXOS*O zh1XuD-scOb+1psB!@i56-bNr#7{|Kn(=ZYx2qge?Dg$8~wXX0hsX}TvHM&gcN$|!0;fRaeC11@Gh)O7Tx-O({U-1;7ultyVYSCR&|SUb8%G8^di!lhhkE-hEpQT@yKL6Ai`Tx4z1DW|+B{Q6 zDdoQ?j!yh zue3IADPK^p;7|RVS(o5iCxv33bMPX3wvIM>E~Zq%7Ai+#Kl;ps3X;?TLzT3_$QM68 zMos0%etc3%8*OcWFoX&QH2Ypw(%SchKO90$Uw{7muS(kZ_!r+jK~>>jEq6s^@bBXP zQ2e|^IM{Q#LLq~9tw30#CFw z$IUOcrjzTm(7#`Lgz&lBNbPeP(;MgC5TnE`Z3Y+Ec}}~;BM|b^P)qkTz>k9OMo zRA~lOaT1MUsXsxL+uSZteHL)?`S5EtbP4X~cNMr%8q$p(8wBq->f0jqv{E0>)aU7_ zZ-dPo4aZ{j?dpYfC<`K~O%HsQN!bf}HpJj~&bblrt3pa5^97DxLr;2~!R~UPWPBE>r4jR+5YBF6lYn@$$her9%$^L;&?OTT60b)URFibp~0X8L5FV^Zuy`rj(?J-AwU9l^k&&Y^&ylK z4N+l*h7k2KFq-TVI@QSU>nRTA%?WmX(HErHA1AQSB@ zW`Xji@2zp#3V6yYfLIC=lK6>1X!LyKxAydsy)w|zVw`>rTk9{73n^|8Y_cX*KOpVT zGFu@&mE!hWStTp{YD|0a9$gP;IEldjTtbnlCv~VF#`f^+i?o`qG7}EWkl!y-hH8C& zVdH^c(P}UcR79C!LVbv@=+i6L$avd>k`xtoPDW)*)K@GcO3$%_v%$QuD9|@*&Vci? z#aNXj#A*My(CT9^-{!OjE@-=TYVI@HaxhMO6go>}dV=&K z>zGDtaoR%x!^Vb+T)Nq8l0C5kDJ^SERL-(z)|Ksr_qN>{4*xxV#Z2ahM46ZlJ8J%C z!h@huN68}Fa`m^C5u*>QW2R*OFyT{UxM|IlktxD`nn(Q0F>fcz_9=x0OfYYvV!304 z=O1D=(G;USKw}#vmZ;LN#Rt9VQ7k+rC@ZnzaxdVQ0#p;QXhZNyP>d2&#@^w;FfLF= z7Z94IV%DmIfLJ!Brc7hmey~0h2Waz;K%pcXbg~8r&^n9zj0p&>p@7O3i#V5wXJWix zmu6T4r&R;`&!{?{ln;k;foxXlNOsv=4VoXvFqLDn@ejh=Wiypj@mOdL#Fmj<=2(Ul zAM_Q?+zVEbu9?XRQXUyd`UsA|QW7g2DL>kP*3+WyL3RB=BL>$Y# zTK9K}AMEsQJn&$Aj@Wecv8(Weq>gwQ$NVrqi04Cn#o1C;O>}ikB;?JdjmhlI7R~b< z=GmJ~t&gBhI4-PN<0)974JtNH0fTZnnQ)1+HdrLltylMz2lFyQUvY*{SCbX!zSKdR z%tT-^qNl`ElVgO|wbcg9C{#npco`f`bIE`m>gg>%ftu|ze(CgFRk?lm1z~$UT6^sm zt{j?cE*DdKd$a=H9Qak6UTBjgdN5t!$;ON3o?<1hd)q5cK#jtETepF)yko%KB@);v zY}a&j50!kcpxDn=Htw1@psJ!iCC49U7n22~h3o;N+3lIk#4tL_-1SL4E7iVu$>H_n zZnHX$_sB1+dE;xNFOJuV6jLe6SGcb*e;~~Q{-Dtq{xfqR#Gt4J1MIKHv5MH;Wh%@z!vDQ< zk;zVCVC_P|XJD#-qEeu<8jg@c0fr=@(f5(wY?%KUYl8M{pBd!53z z`$qrul<~QphQ>gkHq%Ob;U2~0o z^qTPrcx)2zg!r?zfq{+Im!PC@?78V5w`4CY*bT}(3TmpEzr1ttEJgcLu$DWZU~&D) z^+#MD`-43=56UH1bJJYys?ZlXliDEXPwBk)R3WeBxK6GU z(5EbF1$Uwi%CX9-`g?ou^{VQX;XzCnF#@DEJx`5QkG$;R@!OB6_E7Kkm=bkgC(k#4 z6a|YX==Ry#HIi!|vopI!9bwn57rBkw(dMmRNo&6NmsapaF8QLW+%#f{?gc)paQO-L zLOb#CygBGZo2jPy050HUj@#Uzr9Irz+9UqiM>_ik(=HASXKzgm=s*)t4M|3oQocJ9 z%=7V@Xp5_<_SVSHv23ygm>Co}iPB6xWh9`xganGwfWx^G*!lcr9S|6$G&ZC?=uXix z)I&JUdL+=LeBFMh`HDUM=OMXdrtwj~^_^%V7i|HJFK_inIebhp=lAXD$yX22SCf>n zTb+b6?a-Wjy6PAt$J$Z%gNjCg!BAU966KaT=*hdqP+vF*cG~KFX(b(u#kgJS~7< z#da?76Buug|8RHdL|f->o2=)T-k!Rrs%phND}niLhZ8TVxLTf1WlmYxJV!r?I#Hgz zpCP)nj&P+eJdSwy50uah>-1w%Fc=d9UL-i#l2Fednr4U%iB~_R4@@ zcmJqsL(FIu-DKiADW1eMMTP5q34a0dK83z@dWP z*Q4bga&-19uVe>_UkASpuSX>RZxTy(_E8#QT5A@b03>=p42b|7#$B&N{vQrokQ*3> z)&{`g@UQCicsAOK#^ES$HghK%V$9TN;!VG-(Bs89Zk*t>PHrjcJ6Am|v_PGpm;W1=+rd?|n|>3c&Cb(x zl(7+7H(JHY{{VPd+55gXhDrFt+BfXj<_mk#D zy*?z*9}k(bjHfu7^!)SGt@c4Do=bCtA?HL}4T9q~S!Z-#53=)6wW|$gGWyB9s76nD z3j;ZO)ZuFOP3?hm%SC%=ncq@u4<9O$TgvS5MA@N%tOS1@-9eU;{@A?Jdm z+_6Vr$Z~{dE@V5xt_!*4(OHc_+8FHD*;D5)sIw}@&xQq9=I?)zOuG6~a-n&Gr7Us# z(Y($Cwj3&UbUpFMR(D0{+pWFj`8vnee)jte(~8QOw}w+2c6b-3|7F^|sw56d@qc{a zvKnk*EJOmfD&iqurIdj8O4QpHFTVgBm zx5U%LOK@VzTCpiyKnL&D5sVszX?30bspzf? zCEnJ0pSE>k)#eZA@{%p+u97v0=&s^5!gKX_>ZZ9~{gml1OO@ zSR_X}H#}L#2rm?ez1gYBar>}7^CNyS2aK$O*3IfbANk$y|8`V~mx(34J#aiP3I*$_ zdFQ!6vFc|5Al=$;#yjhzNi1eOJcV+5d#G`|?hOKYI5-`v=P(%&Q;q)A2~MD@nnsrE zlB}`>{7EU{SrFEwaV^%ASSjcbX(0-|Ph@CJ1@Yrq6823AlB1y3rF5YK9)7>A4MM!j z8mner%)4b_8SmpND(Ar9#e^f`T*v!bcd)_>YtGMhX^810u*u^0 z0VVzU*6luiAal6g?vc0Qp5J8}*QpRD<2C5W*kO6TWznR?m^t2d>Ran#2kF>;vY;^U zAItv?WeZDGpKT)UZiz(nTJP~6zAxA;Qh#D9QxU*vHnp`yAq$uT@46b=cTn~*(u@Eb z)ZYO%7~nd{GAh9rq1(ZiVDBRbN%ti;`5r3Ye0pl2NV9;D3~x2$hz{^wuTsKkt)>Q3 z2tLxnH74kh_mR8PU|$TpAqapq`53sd5(Bqm;D99KaJjtp~cl0 zdVCtq!_byQ4E-ktHZn0VcwN#g=q$&;c?U7@A`dLJ5Cab+VPIb}02c&d;N%dj=ynWD zJ&b`Troox`NQH$vo`RwGV`zwrq2uOY=spY`dIUpvJOrt;F!0)NjE^efHbCAL41CF_6?_G5|M@SZ1jNSx diff --git a/piet-gpu/shader/gen/path_coarse.dxil b/piet-gpu/shader/gen/path_coarse.dxil index 9fd593ca43dc26f919b022903612dc81fad0890b..b6c9398fd96324a92f6916a55a37f0d0c4376553 100644 GIT binary patch delta 2074 zcmZ8heNYo;8sANJHwjtdhVZeW2+ITlXb1~bg5bp@AfQ$*SafKMHzDFcTh8OIIAtnr zKKL+rp?oYjpdevRJv#j`Kxv!S(Ja9;wzNdBOu6e-P_%cgy#}t_a5vn|?UKZsx&3F~ z_xZiQ-}Ah`=h^L74XUJt6}kf8{~nOp4~|KH%iBG<{Mt~;p`4|~<)*^k-1!}a75krA zk^q1KfPoDN0N6di9LE@jmx=2N+}|Xk0Py3)pLv}eP)H4NW1`~Djwr%L%%ia&F{Vi! z1RcLZ0D$TXJ91RgGO|b#f=um3Do4^nFLq(V0<@q2c#2bF#gKVuRYk*Wx=61F$CHUw@Ln+g8ZJ>1gSLvD769j-4f)R0B);Vb)1TPOS( zcJZZek6piq(-~`7XFJ4?-yXeSAd8#}FQqdYiLMVfwgmajG!F8u1S;C11%f43n!l6r zQiccstHIFDkv{Xr60eaob^!z*SxD{GuJasSJQ;#8ux8XsHHqSjRKP5_dhkxp9gQ)uAdJEHnIcANnCR_%iDsvCB)S@#WwVu)zm z4^ObMQ{GT4M>I(wmUK2v+Le&_09VOiNu{AZe=N4aZb#IBX&^o#D0 zI9Ame630L)Awrtn{UjjC>~vcXs6FEFI_#dJxf_xbtI|voz@X;peYzw(ux#1qr(TXY zf{LRAMJUP?M*LGu%|!gd#F5dk$BxG_y|Ubq%OiSk*EiSo3-F+Y=j8dniQhdXW9>SX zDLbqKqE>82&9b2Fix*njPeaBmw!JI-AV;239g(0NU~S)C%?ZxQm$wIe#b!CHQ`}jb zSZh78XiU%-M4zPZJlW^V#8uC7toR=`die`qH*v| z>S%O^Y3PPPiF=jf-Lp0etDEGKo_X)(zM&7I|h*G98nsOG2U{Q3v#zZIxgB!N?O;oi3-s)yN(>rtbk za!mbvE9eQg5Xb~N7iq^8OfSLr_uhC#^^;`p_Y>4W3=x^Q&fdIq`BdnV0>LEpc})4X z5)zFjQ3n#?EBbRqCG&KL)2$Z{h?fqpsr73R~NCLtySv)yIh>W=ES6Lvwyj zEP`Xur;YWj-)(UZsZ~7Elo$Cwe_w@)z{~)Q#J6^53IlyMtcY>uZxZCicihp3fXCf9 zJWhU1Lbc>)zCNv<6MI`>QRy?Ud2m}n;7ty}uQXVN*aR)o=(VC_Bm7t`%qe#laTf zHrK)%h$1GmBoBp7u7#^w=(tOCy(i8F#`a>58wz5pXa$X<TtN6TdyHPU;G=F3&42m-!CrGS&}gtlmlXL{B{-<$XC93*SiE zoXhxkPr^`EUE1{!X>vY&vJ5NL<+eVXYdM_$M9d{oy2^;slSJu#B&l?y{8fO0pQ8Q? D17*T* delta 2185 zcmYLIe^e8965kEkWJ5ySBnCEsuowtnAuK;M{P3koK#+iA)sEF9}c}V4|9?#ycVeMc0-TBOX zX67?9M>L%pMOLvPlYgvp>$-ujv`?PGC8jXrhFnG&@5RK)rE;J*}xE60JpsCij6pF$^Fp%#jbDW+Ov!1vNcig$<`Yq z-yWL0*!StkR)Jl$vtIpHdCgl@Z*JrCiWNsMB!})mOS($G3@K53*8#+4+U~6V)rOyn zqBTVuG*PWT(7_!e)04OLpZ1MRU)(B4p?d>{ox%+qb6TNj`+D6)SSLV|vH=3K@2})y zFfp%Teqe$^L72dn4@w@z2u81?VE4bz@8eZ+m-lo(^m+X3M)mJ}fBSz2CCh6QdcQ@f z&)z#HpgrCgxO`KSjkNE7-~5aaN1G)U40e0*XIXs4Nc4qf08sJ-0AVsY42rxzdHMVB zidF!Fd6#24^-FxGldt6QXjs?pCN1JH^SnNhzaQ^LfT}@j2ja4W91Di-{LstKja{ay z4S+}!11hGHxCUcZm=0=iDtRy+-AcJ*HiEIuDy|ON=ZtFQ6k)>jYxoogosze#;1{UvWVYv`N zi;KwB^7M;C-u1JU`Zru`LRr^yCs`Wp^pbW|5>SI&k85n>s;~!&L>xf4jXO7GMWViE z)22o8chc>mHn|{MB%*7>9=DMm>95UFB?XN(qoK59)@fC=v}3Ziv8Gdiu1xT-YU7(> z=`%`3vAufbD+WM5Z!c<9u1J67FLuX%Zem|Dtq*pdF=|?hD{ztHe z?3R~qncJf=jHqxm=XvONg| zBjj{Y;i_sg=W5I%5Pr1Eb!rj6Nv5Dm^XDx)9e#xW>>dA&PxMyn57Li*9B0Qnh7BQq2HYQOC>lkZYQ8n-zDN18Cw}fI@@#3Qbtsw2=kw zO{36h{2I=!#xHmd>9s~f@MN4-V@y0hirO=MU%UsZt0RQSFeTO*t>Uudn}hWbr_hs~ zUNOkjg~awr;GqIz&R&vKq{!Baz6>~^%hOD3D$LV(pbAkALS9w8U^bv8lT&qF$0yze zdA!#b1M0A1+T==LP1uhCnkFjr1XVj4Y6pWmkb8BR!R1sOg^CKXcW&w{aHDnJ!FNzP z$ek1x$@7!LQ}wM3cml0|5B)0G)gG2uBN*f=c>&BF9FoNyDsXNg)OXv}H>S<<(Litu zUg2R%{w{5DMgg6VT#zO_XbsoE{es{L2rd=(2OJA6U=(4M^t8cIE#$_&>XTJpbk@u{eEc%2Rg$AOJS- z8Ug^C2B_(L>Wcx29~ZBC-~$7|cB z%}LO{p6^|4rt(Aye&(gn) z3T0j88%?RJ*~OeLp0O09X3k6ME^q&QIsY1Y1O@K1z>z3;3@b*}nxCT_4?~2vP57nH zu=%l@HX-j2E33Ul)Ew=UnDmtVX}c@lO# zR+W8kH0$VS`n}Nx)2aB1A67n3oPEgZm!{{5B9Z^pg=5ei20I4@0GE1!0F$*gqCUQh zU)v4i;>OMR3&@ecUBhyF9erm7NR^uYJnW(^gcxNnWrzl(M)f@A7l+E327t4)uWZq{TR!w z15gfhD9FD-{L*T-+{z+9s&oxJiF0>a#t zx08MlhXN(<^#(2ggQ^2qIJ=m19#+D^g-lN&2Jl&`lP<5&j$U4=Q3v3B_~3ezimVY0 z37AYkj3xEz?UczTd=h%8&n;4T_LFLSP9PysScQ^o z1{|!b@3CLzMaZKOC>V$W9ur6hFafAlIV6K*qP7(Zu8vpl(Y6U1qky#4_wwN%Npf|Vb{in=W87yK6LFfXUC}@z| zt<*>B)F6YJSSS+PpXWCi3*!LC6}DKeoc`jI_gfavVFe1QJ^rYLS&g6oN2&PEAy{Ff zPo}pw=mlWeaz(8pIOqzk*06JXe8)?V^dxYsEPpjh!WfCyGypvB+sU%k@ z0i(-L(ZotES2&8=en4B+FBp!lV|5YZRo$F4>{WRtCj@zuJEy`4-mP~waw{$uRw=oJ z1cZa-Vm8(RBsvLF1!Bu%z#gUGwrUj=cA%9zt@fV?8BLj2=T_KV9@|}>H618r3Mam0 z)ovS&<%_!aW>51;g;VGZzi7BTslR+J&0i?>@n_BOKYUugbfA14%}*$z`OE6yR3K0O z_CS;bl=gdXq>Lz^CUl_e<&yjORWe|j%RGY0PV!RW5+yAa>a1xYn~drYq&@-my`sNQ zPx9W_IDAcG-u!Fsohuv^MU*OCz`jExyeGk$a;gJW*i1;q*eqt(CzWjp2KwU4(r$*o|K$G8u*-M8W_j=kj7BX*82p|RP9 z8neFVUSDo;W^Xq0(hj=M1V-80Yndz)9T+Zf*5IZYkN+z4+AOH+qdcCCO@uNX7CJRr zSw^ z?kH!sJW=9Jg`8Bcy}-I{iR0j{-fce66piO(>dUC}cI_dwzQ6jthZGwD5SQ`cZsaGs zySVBPWI0F9{ryQOZ~)H6(gnT;UErvjlZ#6vQoz?@TxLzPx|lUPj0S7z&yf30gQd=* zz+TcMtWm{amV8*RT4y&~RK$Hn#N~UYEUK z>m$dx{pvIs>1SEn{TBK4Ke2M?wIz&lEM6*@y}{QTJWM$W_snBhcKX@=!^2CY0rcUn ztex>n*};G*LC1B<)v56TGmyu%*Y~hC(U9y`YYSvR@aD(fU<>VcvEi63d_rhA7B=By zI410J%|X4xjd*Nowt_P$Tr97h_W|m?Z3vHjPEyjP@1&6Rq%cYT zjapCVw}>lo!IfC9l+ISXo`y=_YQ%#|Qeua>?Y(?I$ciclL+Ywd%qTU)PF7NpB}-aO z9P#MyyELz^cGEW0`}=V`Hh@1VRLE;TVw80`n8Gb!UmBDBEMVWxJ%5K|0qd#iiZEJq z|1B(pCmfaAKVeYE$iWG5d!;CYd2jL93TO}N9j{tcS8L5GOF>N=ly20*B59*#(^kKJ z9HOyUt-1X6ouUixUAZ%3^`uQtgfUc+--sCdf`#f@>m9z{4lftSCPIv|YM3>Jc#z^U z*mekw4VcL`rvqxMuU^_|)>nh7q0-atX9+WVeZ<1{&)mi-{ zoNrQ1HIuB$UA zxr{%Gr>4~PKDBX-P5u9|41$lggbaveH|1*;W4&|w`-=6>Pd6xl-+rY1zvHm{PCl<2)Nyt#v3ixFO)>=g+`1z2| ztplmE;dT8Bq4=WsCMD^x`K&9!_W7Nto^$;&efDU8$a!A#&F{yr`C4sL92rBRP)cafw7fR;<`Jg0mO2 z`9VC&Qr(Y(_e{&kDAoNA&HcB__V+nus6=zNZ3l|yh0ka^b8C!DAw|ju>2 z4&hFRXZSzN-*q8O{P69Php(1D>`z%5HG}o_tq0l^ljz99Db2%yf&HLXmM{ieOm~Ud z{Ui(==a9_#}Bye(IWW^{v5%Jv!cS*;lq{cuWzvTSrbXw`-N( zulVvP%1U5mUaiF0i~;+?Dv!c~k5lpvFb5}Z?GN2b9)2J`kYF3}Bqk`V80Rkv2MmyL zGF-;QigC^~wRkd|DM1g$q>HD8>1@hM`s5(6*D9EGgymwA2Rms)R^CkE2tEgME zcRtfeRY>ed^QawM;jWwlRT&j?texi>dB2;!@!i*R z{rKCAr+m?ywnbGQj|*hKesDN`bkgn*i4YF15MRYQ5Xec_>~Era&YUmHF{A@&{?W0Q zAajXJt2in}rQ+WEJ}?vv6~xd&fWSu9`M}Qe_S(w0>?)7A+2wl^Y+Kvn48Hq@GJJ9m z$B%0z7E}Y7i=+Z=r0bnW6-0%Oh&~Zrd3Gh{NUShrRM^K%Jr7OfvWj+TBYyy&S4Y_5!d{prV^>S>f<_`Q;*`iSEBY ziBz_pbrVL@oq2_Z7(bl|+fuDJxv>p-b1FlLA>#-it5bdB*rTWcLFT-v6ovc$W`#tg zx1C0Lmy{8AxILax=m0>U;}dWVNS*9{9o-)1!j&{(L=hPfqY^2pKDjQD>XT6@TdRCj zYbxAbx%<1M%X+<%{qsxqKRR7K)RYA5(&Ub5J_ zbyj&YWm->rDmM^GbZwt?mnKKr(`YiN(@ajar_G#KwnxUS@r}3dj>N6pixvZo4NvbeTFuc827}re zU~hzcpOA{=4$E6MGx+DX`68Ep(iqf2Ib{_}?>AaLeHD^S--{-vFG-x$fsk&JI8I=r z{ALHL@^|SGmbs>A8yx3pzP0p=e;D82S-R}U2)f0aaBzO{s=k4*3ypWzDrA3k^!~-Z zi%4Y?XD2DC@Ls*sa zhu!PGTP*{uvc$+lfYt7>IZ-5mxByvFB6E#CTW_g0XByW8J5Aozttq>@MtE^he>^S3 zHV{Y@e9X2~IKA`wi>Z@qde{Hg-SxLq{ZF?Wc=`_&`cReUuoyF6U%%OO^=`gmx&8xl z?r8FH9A7x=X4ic%&}K~!!rm3-J;7>v!)cczI`dtT;;#ALyx>tn=5aQ`K5hy5{*a4%f_s4GQ<*l2up5=5 z+-j9HaeAA(qHG<(W28P29ca|_B8PsVI)8*S9bhXIj3jWS^O*Ghlu$I z%)U}!V;8CmnX&E@75rk?N^gpkEeAmi( z>YXaVxu5@X+9YETr!w9{gz5AhrpUOf>g-gA^u52^n=_FqNchU0ZIfg}wPe!`HJSb= zO_~(MB3oUmkx;q|Y0Q1%pXuv7 delta 5206 zcma)Adsq|Kwx35bgfM|55EuxU0P+S+00~+TCcNcgqm77Q>kvdy5F=Q{XC4DF2SHGf zV$~WIEh^R^YO%F#0*J8%r5?3lwFaeL?BTX(ZHuk9o*4+9^PTT}-?{t|=C{_~zqMzt zwbx#2UZ#zwc_wDZClFU;SKB}U0Lt4uS2@n+9_oLhExp#7{At){ z_$vIT#O(PQHa-C0038@d0D!^(Ih;+-u#)ol1oL;kC;zF4;I0A;O`s7ci8o^akW-&WOHc!h5`PZBbsF#KcRO$13uNS*GufJA9Z<+eO0XDq#l?5| z%y4){ZKdb8xtP!4kJT|BM|Vg*iImK#6Q38&PJV0kT2*m@FkxmwSeQ@|FO%^hag(sx zaxsrVh=l|T#u`A+VyxGUS~Usp^Y80H)|VPI23=i;tdLjM^nKJ9P}AC5Q+`mcyB>1D?b9N?C>e}LSC=y^f_!0J6YBK2&y97nESe8Vd#%T4hD zX(*9`NQ19`V1{qJ4&vcPp>Mnvf)S>hhi{7Gc`qcO zNd_@@)H&GohdMojo2ASxcD+Za(Gt%*YO4QVPNw>$MTr}ISm>Kr?IFKO!fC%`Rb8`I8;ItgGf7?%xS1|S!9Vd8$d`S#0n(~x#!qT-G%Xh z+d^lNQbT#+{RPIv2wI4k+}d{!!qu1%VCb}WT*I_ZU^=y{Ax}v6P3*L`GJQ0DjC;%x zl(XhfQRZ`8KWAGh%3dx8^E7gG9aAg^gDX$6Nu)4Rps?~31u(bPcvNEOTB@;x!=5U)yE-MiBce2MTO$?y_<+2+Z`ZEP_83Fx%^{=F+~W0>%S<<12IGn@NvBJJS<~ zE$ZH+x@nX2iz^-MRF4It>pf&!C#jqz<-b$Wnd#QhIL%Ia*U;^-oc3&zH1~t8Q{fe~ zNMPg)TqVP6nP@az7HyB&D%;gUgKWpdseiHFeIx~$fweOJv^pVtwOC4?eK;cd&xBe1 zKw`2hWQ5qWo#Lm%9%xp#VVF+zMb6#Bck`~83WY}hAU{4>H9;0!>fW46kr{Z`NV1X< z!W$%(BEig1msDCZ$|iG#)YTOmtPv(t%sOng(~1dZS`Ed_O@!66;JD~CL5qzi^D#$enI@=5P0k z*Wnh&9kilI($%*crcTJOL2J(mvGRtiOgDEiF7@(T_{L^=hf97(F&|r`)w}6`4A{@WsEl&kukB1rJh=K=I1Zd z{OyL0&9uEXH0sn5J+xF5$QBTA zw4v5 zxWJaGZ5i3T7fk}jn0OxB7L|^(HlSmo^==FQ<^Lz5HZW;1>_g1RD3&g~Exjsb*3XoW z%gG_lq8Cu5%l}iPN_<;=kNI<7B)$aySmpFgoNJ=L$QGJcOttt@_9wZ@iDqa9uf6&sb(8PO&r6d8_( zi~bE5XMB5()cnvHQM8A+bY9IW2`?3Se0;4Y$@%fKm`~cQmSwh;5XMj8K;kT%6sMes8SD1aho>K=wL=6GQK>%zsD!6C?}|KEPL~B)LQ47`g1mh(QLE zD^+bJV#c|xeWsYA_Lx_T9?2wf%Z-ssuqnXEs>Y(p2vvi>Ns7B>TE1nh`#EM^ROD~Y zW_MYPNZC(C^0JYRnlEVU{kt>ob!T2X=ndo|_sn+YW3WW%+FK)%ltoE`%F2LZ-_@`` z!Xg9Ryl?kCLxx-ufB#nI{pQ}ULDTULyESmroDB()mivrNWs*(pkxLxRYIj%3l(ioQ zYOc{X`wwS+H=Oy^K_4LZdhe9A(}X10b+pDL+1f6dS+)r%mX%LgyRiGamW+N${DbG2 z54tmwR!+$;UH`5%!zOt(B-!yu68YBNzA3J=U6OFumwmrV)IUq2)MZMbc%fu!(v{s0 zIx^UE;zueLjNIz20Zj#;O)XvWh7pdkj4+Nx%{f*UbwQwlX9-LMuYY%rR3Hle!{F0u z+jy%ah#})>v)cURTA-bF*U#boK_Azp!_uhBWq!N73r}p4uh;|kjW4J|`Rwi+sM(|j z#V>mS5_M7tc*ZpboZKW5JfS4^^A z5w&0l&m6o`zoEmo$ybE!4&MDR%(s7t0f&6fg$@-kh#4(i`t!Q{oqWg#BAhf$1!Hk} zvC!Wt)=@uaEAW|sveCwG0cdYN(erg8h*+u@0R-M{Q&Y9J*DRf{s3jS7Chtq*(qkAz ztm(E*C}OJqsaWmrNh%tJC^843p~JQWdWkV#YL+_3DQy2iRE6Fd2@NLBNNdhRB|BM+ zC@FQ6gNb~y6lO$LG4%`e zA@eU#aO24dSQR?|5(RgJG*a5#Ve>7E!b#k7lEN~MDAXF@Q^F{9MD_j`KN=5ZZ-~|X zTy++bt@6BSM0G!u*E59k)Ylp>&%(s9RL2zfLCNY|f z-vmIz#DK}(si3)M{z3A@z7FQPsM)OAy zGHsd~6tnmI&EERFy0u{QhV&n!$x?B`urkXl4vvO`WbP%!g21=GYKA^Rw};KGDd|v!C?pSW zRJ7B|MEzo6snB7)ky-NJAMc5f&Gn!9`KZ0>?kqMVxiV}t&|!OIM2r!~Jb+x4DrX7! zw#9x?ojJZFl=u90|B~&umxwR)Sq^7~IR^vDoN~JTy0T;~zn}nephB&C;n#|Q*5h3) zOOclBLVRw?L`s~x2yd*tcDq2EX<21Y3-^By+s~Wo^2LhM%ou*ZIQG*odL? z3;uq7;~IMb@!cn<`*eqNH|&Ml7fCr1-ZCIVrk=9>G;7WG3Udxr)aG8Ky=^Zqh8E`> zrsMR(_Q1BLU)T}a0W_Uj=9G+jkts^!H1nHhCxsU)<_%Px83;W+c{0m;(=YUxb?07N zd-6-p3d-8*Q&ZOVe^Xqu7|+$M9c)~Ca`BnYdfxOD^_c(%_io-cd-gjy^ZuICHV98! zefG5GW4s#&$^6*z$KHqz-C6u{Pm%p2LC_b4PA=i*u-B#O@}Z<97P|SCC0C?0QCp@^ z4>i6$xTNl#Cjr$MMphb@pO-(n?chHC+7w6p{D;#vHH4qan~90DsWxa$_qJVUw}h3S z{f(PE@ukBS()OH7cS?1Nvj>l{X=_cE3KY#c?oo|}leFMI<#0gGBi^kg+a23!o8=G4 z{!u%P4BGf<8l4aJ#z73#^C1pQX8gaX2uWWzcpooiGo#RA{#34iP&ehLOpb zy7c;!Zs89iH6mH`7+umJCYXsCCQEvQnGGa?xM+i!P8Tuaj~Qa(KgReLX-v%Sk2lZv z`##V6exCQ=EGfX5BT4QK+Sw)uEH9ljx>6D;@j=$t-;rhy?_4H;jX9Km$U%z zKn2by0LW?p(@9t{3w^gI{f(6Xz<=!m_7EW&*rvjXBE#nfP2b4)FrKg)ZgG2BIdzm1 z6-ixW0|WQWV0k|%=!5m=X#hOqHM)*~5oI=5M(c20uT>b)mW?q|=OH{EU_|+iF+%4# z2xAMXnPhr5P(&)n^q&Y;TWlg0yHde>xPGQo2gxMJF1 zes87FM|Zz_@(WA1`RfGpgeUf6>GSn}diOfL^-lZT7A#IItl#;;b>r&7`f7*v;I&Ux zuPesq45&T>8#AIuslWfaQ z>8uiGCF+t>;_&bZ?I-|cYAp!y_h*Cl%Rzd0SIx=xBw~JZSAC4sCxefRay*EL3{4NF zQ~!mtF4ao|Woj${7a?W2t&V3WNnj#4c^bM|b6tK1?cFhlUzQ7ROQC0v5ClNwbVgZ+ z%>ZSQg9Rk2oX)$;q{@i1OKH$T5~Iz0SfU(j`AQ{lv8-FxhkV(q z#}TtADBTTeJ|kkI$b=nJR{}g979@ovHqeg}0ZGB;F^1T`2iEcj^4P0WTh{Bi7l~Oy z=a_J}TEM(?`}f5WlX^ToB1C1|cCbrewHGR&-BUtTrX@vgNcDs&%X33_2yV-v4Oxa( zG0ZP5YLRFtcg`RJnl z!5YQ0D;%A>Qcl^huM=sy9d{|=Woxg^r+y+wBehhSgZh~Z9!>B1`>oU#ITmqKHW&2^ z7ivD1o2Y|%WFqn$W$&bZ<-$!l{{YM*Cn7ymxj_BKEj8!-dYDFDjto+J&rlE9XLts_ HC;^(+@?;a4n4CL{#>{vC-KtMaDktP1Nn+j_3Ek z=Q)1Q9h_pnZLjo&I*%(0NC$UHc_wA9oz2ib6w9CD|)sf z&HYhdDA4<;7yuFI!9^Vac?PJRK@(Y&cOKW)Z8QL9=Ppx!VRS2htZMqE@$rx3=`-nZ z!!;W{a)KBMQhKNNp;x^ER&HWSJMImWXdHYD(R}()Nnf~%T;dw2iq}LL_^K?OgNd+6 zq1s>dxiCnbL3M1Uot_TH>wQX`r&J@+k7-24>AxSpv(U%1jLIO$ z63QxnelaU??km3r--6-?X^*}U6Vq?&ZS0#h?)ER-p0wA+kWtN(=;G*+HS0sa+&d(!#s#E1YTdZ&Z%9c zS=kdTf$-X5QAny*_3v`AI^6>$D(Ejp%PC$JO)EN#w9ZR61`PaAG%tI#X*&mr2>&Gm ztQ*oPt;nr%r3V71U|cbD90c$=*wlD=R_-qd43~tk zCRIQ4pTKNEK(Q7C6!}Js44$m=4vv8%53tbloEXPypKO*6z#XB+8%FT6tqg%wufzd%(RCF6gv2;o@XxO?4RbbyR=0R70mOa1` zpcjU8eYIno<=?)X{7$rr^R{S_dwPVK65~1hGa(e_izna!?5(R!SjqItCqpQHkyCWE zOWVIBUDc{icG*=~b~R5&9FTMYVJwCmc)~*==qkY29t05VebLnPc+p2le2If-l{$Jd zzlA2|VXHpB*UqO{_A`gbuQA&=eC%cfevLOB!pJ!~Zcai7C|SC|{<&eXW@Aa%f-q-i z>pr|{!UE7PSnT&q)>&fzMhjb?F$n?Aal+!rnH; diff --git a/piet-gpu/shader/gen/pathtag_root.dxil b/piet-gpu/shader/gen/pathtag_root.dxil index 77f12e6db0ec75e4d05be227c033a9793a2cc1a0..48584bd22a13f8f01c186da6383ccc3ba77385b3 100644 GIT binary patch delta 1497 zcmZ9Me@qis9Khe*wbvhB+f$%kSH)xe&I@04m5MQLQ%##gA0F5$0nE=mxV2CBU!eXZSNF~+xz32d_M1c z-+S--UfxW~Y>L6gS*+lH1A$zB$#7b+wxHqg_zh1~9&_FOxGQZ@la1T_YOfjq9>~E( z6aa|ml828UZud=Z>bI})s8syeup4y z-Uxc%0c0k5btQv7 zOFegEub_$adjJhO>3DhTr{SCd6N0C8UAy;f^M=%hbH{d^nJm9^_S=&}wWZ}j#p3H#Z=~g=GCu{Uuf)3j$2-1Qb?tb^ zr6Sdd?8<|g4X(yL`x|z;8n>anTEn?l&KTb-Pn_7^W&CdPB*6NQE~r6Yu$w}Ry2p&! zFF5QU?Zd`Gf_yA5gm3DQP9+&v0WUxG*2Zw%MEeqLj$24&bGt~eDiC>>qJh%O_Ta|2 z`uA0c10P7nl`xsS?D<1>x7q@Vzy3LZpaMPdVnjY~`I_a^x+xm? zVNLFk;42A~BupX(6`)0< ztfT}dc8|oNX1qvPvzTWl$v%{dWh+^xf(Nt`cjK#@Et!4!bDqo}hQ_SRX1o$0Su+yM zp;^*Q)*PC^I}m(6^f6IvO-4HLt3w-8Tk|`u{X?Yg!IY%e=EJ6T6Ji))Q!v<5ViXWY za>xXr<_R$)A~n{sBpx<4HOiVGxc@_sVUKdJqm82AWk4*t&&L31t+O9*&=7ExNrBXB=U z$v6)0o6u(~7#|KJN%)AgdZ$_14}uR%aGrr%XL> zFLA~>IIVG@M3o_!ykfPW;GBc14rK|oyCStYq2(;!>hw>im3bw7dPOhrxz!AeQrLM4 z7<$1;tf7z6UAW@eJU5$)^`0r#uTp{>v#jA9evBfRFCe@h_;=-fRE> delta 1493 zcmYL}e^3-v9Khf1-tPU{MdyPA8+RKdGCGS z_xs&9FP3^cH8I~|vHme|xADuD-}`*#z9>-N-*O!R0MwmMJR)98JU{gYe)+@5BdTkY zX}-929Xt1?WL6$s!c>pN%4f!BZ`_*sYJ9d=v?W<@JfhE8J8bM5G~}EaUKz&0z(IK} zx^lcH9+RCg$qa#1*}F5Fass5x5!i?A;I~QvNFk)4m5c7%5Uju2!NeI{?P>1pw<)m3 zA3msIfTY!Yl1NgwoRA{6g25C*0vA)ChVO|u7A`v6l2^e4Mnb5PeC~LExW)qF=YRDf zD9>o}n{7l1N5GB2dE0q+61j|WuLgO_t4SBaK}5x3P){Jqg4RP_l6*zbj7#u8x)*#F z7!ZX`cnG*Caw~`riUmM+2~Y(tm{}VVL}CP#0c9vzMAtEWoT6$3ASQ}jVGm$>$t}*H z1v_GxCd&jUk*6>Iw1w>)Sg0HiAIAwtv4Jz}|FMu!;m7g}@nytL~!H`MbF~xl&vl=_Gfpe5{_R=YKB=STq z)O_59DC*t9HUB?|K)zP?eKwD=cE--?>qi~%mC`oVEWM3 z2Ln5oYO;lk$9!fC)>%|Pp?0w<@2HA)s^}~`vMuyrqG4K{rSgqx=?X2ax1)}S{)Ibg zy~;nDMOS9gtL>gG4}Bv-uU55>n&~PtZLoX)^Gzz+pz0j8)75r*jgu}HTE@s>%WQ`L mRso+!kSaWq&k~RQ`5Cz_R-ZT&y{6wC9;^nQC?-rVe*Xdj6YOFD diff --git a/piet-gpu/shader/gen/tile_alloc.dxil b/piet-gpu/shader/gen/tile_alloc.dxil index d69db16005e41c4d04ccba4f03161ff131490b28..1a97d8203297cd715fe74a88b950c7f381d1942e 100644 GIT binary patch delta 1254 zcmZ8hZ%i9y7=Q12yVSv|ap%ONF6bNjfC7X(}nbHZ44M@jyL0oK0 zTib15x)l`!<26%sjwpsr`o(ai1W}ljh3KY)Fs%fWWt;O0Ci~!Zb}sSx^uE9M_j{k` z_dZYFU;$THgO-n9USJYH?&Hq=Tc)CuEKtuyCSV^ht?L<*M{C%X&q}#TUQ+(Rb?!aFO8&M zymX~Z|FzZ^ct7y2k|+wWtTdEd&J;S2hc_Jz*+o`pFsC$34l<;_TUelTG<)!s_QaCG z$;L1K7@8dH-sc-uhfkj0J@)2hziRi`n7r}z#e0XoU*C@2XE&~|Z*NT?s_9w$BiNP% zF*zX5@Er+Z1Js>WdQFA_03G9<=b&ORVD^2(=st6#`9fX>wo~LQR@jpaeK16zR{8@1=vTm)eA-mCC8G#RwO_)9VKde=_cYEMq46E3Q^KAB$?keE4E9(?6o20YL1PEa0 zO$(E>K=rv-T|l3A(|BP|&dr<)I*T;=t+LWYi@<8&R4^XUv?PZd@Q7JYi( ziq`@0W(%W4LP!Wkt3~%yH8da)!SCty>L6fWge@%0qO$~A1dam#FGPxs6IKS@y|KDv zP640ZydFp%>?Jk&m~1QKn3{;2Y8Ab_DL7Q`;p>W@)@Amp8$5B@mO2G)o8JTm0w?=# z2-`F2@Wj|#T=E>od~we0K+1ck9$lT;uco`|rPbyE87x6$czGmL>VHdCBoOtI{>!ow zn?HRZFiMBwK**lvt6}mRN3?igo?Kz7avG*H!OB8ju`R z#LtfT6GR`aZWzo*48mJLx#@hqjk6jlZ%AlZ))tx9;VG7wlIbR&ob{V2I zS`bEE&tU!+cP)>FAG>0D%uujXpY7jb+_cT x>RfonESa%NW(v;~{owCcq~$J{mwb{C*_C=)O*%`Q4akwH_TU*a{{cOYe*ktthPD6z delta 1237 zcmYLJZA?>F7(Ttdy}hLuxNT{@m9JYIg;jwI17vKnen15gQv$;#C}}&3GeWIE2y7(1 z{h*Y{lA2o<+?=U`ncELjMFuRAwSyo!iG^(F7Gl6K^K|5mKb}gcc`tOci30$bxFi1neb?@}?cQED7H&|F#Pz;X zU%9dBc%={k6X1bG2!PlcuymYNP%N(1g`X=h03MEBMt^`fLZ((IXa3|OQaU68s-qyz zGbab)vjB7x0Gwd@g*A?3WQJ6t7QbkS1dCpeL1rNc=owxpfeXS))Hr5VT)7PT z<-y5?8i7l7rYpD8+SY#Vf~^j^A&_6Ybej0`^p)j6_?FK5vY=0jh&-zo^bZ)F^mnW7Vr&dE|~L|_@DTk9JZwFH8L*316Q>Cf8_ZTz+vw1 zDaV)v{B6e^@BMIZTP&Bz&1FNHmLsx(HX^?yMp{p8QQ&(AZR~4Wai6N}TkwFrU42@I zqN;CCwyxO+nP*mj8pnuFR(RG_J(duw?;Cy->fhu15!v^7=l`M8dK6SkWxUqp5KTE70uYh z`(cfEEyN6n6`UY*Tb%B|#Ua<~AzMZe5YnCLd2uCIO5+QHdJa66kaxX^pMO!h5?NA9 zC`&Y!^||rB=Ll?3_BvK_;N*ZA3d z$u>SzVUabPW&cvu)Iu2;hDIk=nd1^!)03Gp3*XLMPE8?@{3H9CsCL30Wul)$D)sX8Z?N Cy@mJy diff --git a/piet-gpu/shader/gen/transform_leaf.dxil b/piet-gpu/shader/gen/transform_leaf.dxil index 32ec39935ac137a362c71d89ded97dd979d0760c..915248b0d34164c12e496d8b1f9d41e43a5c5e9c 100644 GIT binary patch delta 1694 zcmZ9NYfuwc7>0K@%Vr@PNJt2Q*s=lTl0+r}43`$%4FLsjRAG!2+X$$XN}vc9r_8Vk z*HRoKL5g-%uqqVM27(oo>14S`2C+g*Ep;r4qE%X3hWbNiI-Twc_{05mp6@;L&N<(E z&e;}Ki%P6F=n{c9zAmu}6`VX;V2<7LQhxM;+b?a;V!Rz`s0xB}jDHzylt-QKin>;G z^bF?F3b^MsgLpnd_4)xGdu?)FM!5){{Z);-YshgGB!b}+I6hHbX^@`rTLDd3F)_wzA_kCt~G|aXrD2*CVseFOIJ*qB~!2bIIsPb zwpK6yl@;4ds~+)<0MBmk2eB-4#Tt&=^9k;qTuwVKDINkqLV$&q*sJ1E)9xra zkcV=m8Uhf&oQ}7OXfJ4&bQ4NN6obi_0Dx)A!u_! zIZ!iE354$TH7QuI%r$aKYdsP} z<}>mgAVGr7cMb6GXM^DBNBkK;RABFyugHQto1!Z#HV)AX)rHOOv4-hF&nOM|hU|w{ zTdFXm+yo(md^&}qRTTnw=|EIgCEV9v=4{GzkCI;&ft|+nyi;g3zZREK3X^Uw1TxT0 z$yy;M6VO_q>Pcne7T9i?Sv?de)#4TtkoVO6o71S#ts=8ANxh$d?*WaGlCArMPJiSo z@r7}ULVLk>m+I{8Kx$1n_Gy=dwu8cAj(UTd$;*!SBx=$?r)Au8m~aDlx_87(m@4#H z8Y7-x6s?P*E}Rh;mu!DBy!nZKBB^kYpGhDN%uSdcqIWBlm-sTP59Rlxp+K@vv3!`t z_5ANyTvf_qd8aX#KhE-+GWjD`uI2y8V$#?05`VCBcME@##f&!jFpEz9Z3e3|vhV=g zbtbaU&1LoeoR@}lWI;2+M7GV(HG)T4j|kDSdg6XI46(r(r0qJM>K4>Y!lnG@xFg*QX{iSQw;WF?Hi2+JnMb%uCnT zvFdzK1dDZ9`i^ACFo*|hS?tKur>X^iiMo^Q9`({F<~=iK zgxk-A(4D9wLg~;c*U5J#UvSov9Hvu--)$Ft{I(wp?|BeHUqr7IDX&djl2pCG=Xi^4fq>%bkNNW~zzpOmoaTAej9h{SOe> BKSuxn delta 1610 zcmZ9MeM}Q)0LJgG*W;+AwzP%PqTErwEY-BimmA`h9yOqB%~rK2ZU-nb=STerb4ui{ zNFfzM<&d#0(}Fmil2HX07?^Q^cz*AGu!GmH|)DSCVO|aw&GGmZ|mdpDkCZnPyH(AQR0C!OV+LF{MlBtx66(i zA(*0`I9OWZjnsJk#%q4L#7BFgfbUCz+D}10g#au*Z*8Rx8O9RO4LOvX&Q!wiPd@pz zY?IXr=o&2rwMAtpuamlPTLD_ekB25R$6=-`0Uf(ggnwR*c|2`Nl-0fii<<`|pfyC> zV&lyn?y*HmnEV98Zh07;Ck$@g@v8^APG3Jac)9e)2aUFM4{Yi2=vm+0-r%vT>}hd0 z1f692#8_nDF|RfCj$fU&TsqY|eltHb6X3rtLb?meJ#RVnF%#wIYAq0rX~-Sa?ky(d zJyWp{z)<9rJH>wGgMA2@G*Oi zab?+n!<`!nItb*oVp>R2#k4cZGlDTegqY^EiI{*z)h%dE`DXAY+56lJ(zS10v zMmhBjK~vp-rKt=POhj-tPxN1&fo+|yjT5bRRZoXeGm_3gKPaWZh~{+KVjLcPiY)Vq z3!Jwpq|3zKnV(S^ZI^L`cmVuOE_DE9K!3K$4*J0v2`|#fKr$>+16rVyH?cOb!%A_6 zQNKHlmK0l7N#y{{(Rm>ogH@9eHYMcuxM+n41}r9j5RZI;ja82y29m6i)Gz=Na_N~5 zCadk}fua-~FkBUD%mg5ZMR}DA4>2mprO8L>;&RY7mw7g1cSbc8Eo2^l=T+uRYqSG5 z)$o}aXmBsR0tvoNuYiNg=+ygc!-5kPH&D^@paOY%K2ADSlBXHH(rETIEGofb^zvsK zX5b16XW6hB^VYDblJB&mSB|Ip8Z7T)sWJL{Kvslw=6tqnvIGQAOo!c_3Evgy$D=7o z)|`IWx@MSr>T&r(mpgB}WXsGU*LB}AY`!9~ayx_!g>ktwXRMOLFHZwqo%IWOQIua8 znu4EFf*llb0Al$N(fZ~;)UPh?tc z;>$aPbz9ALaE2L^jSymDUiRxMzP!cgA9mjCo-@bXvu5lP-kh6qU~1%zhvdy1*HBm? zfpkg53P&O-cHvth`PpOjeP|M3Q}4qs3P{vn5-@&0J`xF5c2HFPHv*?t;M0Ol*N+RB za}wX81v`pb>Te0S_B4J`z@GXac=W8Luo!lb$SP-;iaLqtxibZMYOdicyTwN?Hai2!5N;T+ZJ<23-8{4T|S&WkdRI?@5y`g4?mb|cfNxp1M zA0KW9>I@`l$l;OFFA^|s4VfpPU%yzuZ4S~a;H2IxptqRx+92nUJRw!Uv)f3+Djw|# z%UQdMNtno8UbE#xkvEvj-|Up;7fCU<2gufjLl&JEOBMEssOAFu&+6j;q(Q)^c~E>8A`PkDXG zy^YCV?@?SB@k#15Dg3*~j4bJiHE>^|eHW7Z29wi59zj)iJ;Zu`%}fCV3vVg^{R>!d BFXsRN diff --git a/piet-gpu/shader/gen/transform_reduce.dxil b/piet-gpu/shader/gen/transform_reduce.dxil index 63df381194f6336e5c7fa13118ab5248ff4ef02a..5bd59c0abba2a99dff2fa37e30bd43212e60370e 100644 GIT binary patch delta 1346 zcmZ9MZA=?w0LPzu>Gj%O*_B?yl@@z+g@G$uI~icy>Xj}PUy7Zp45l9rC`c5RWg3Ep zq}Niq?giwUt(|;W|Sl#ia7j4z*Z)<>CIE`^9x%ox3xBPh13}F79@1A0s`k z)f^whXNL&sAds|#A*K!Ezfl2cvyk*Y4S=<8MrwB}(B+?vjKpL~8ej#A4lvB3-&acy zsWfmZNNlnaT~?s>i3ga&mr&?bA)frAK40$-V@HRIIIE+L(QN{S_)&JK#IVrjO%^Lq zmk_b!_o)3cbg=DLKRn0JeSZG?=1Vs}ven$Ql^5WrgV!cf{WA+AXD93uOQ6uzvAw;% zwY}-ce)!WY)hBbBgLK`;Z~myRQ~eTRimm?NbqC7s*K94XDJw6C-V!E&vkNr0ClN1)e(kSY`pSYV!^_rC2=xk#+p z*uVh>$7a+OO^DF7SWwh?eUI7im5U^_<{BI4bXlp`px7$Yu-y2#+&i#{-#ujHlx#FE75 zL0VLXSJjv1gIPvIn^-`ih&1PwYclVC+K%%t*&rK10yNHF?sf*yKyd^nFX(ExL#IwW z#w+9)LeKBSE$cg1wdOTD9h+Z!NXDdC{~t%W6ps%>MX+|5)$vNA1j>#_A-Mx_B+OO+ zTERA&YtVb^5nU9EQ0^dn4AdCseJ`Su*j`@7?$YOXXG3H4qsKzHUlwpja(Qoc;mKHBr5{q%I7FV$*b|d9fT8n#(ZftIiRpNnRS5`aneU-amHktb%Ocm zpv_hDQ_nl?=z%ykDT}`xj2n=BXu*_^sz#&UxZ#K{0_dhCpYv4#EhX*`(mb;Kh~p*a z(|F^aZz*GMpG{-Iscz~Z8ubruznGak*zJ5n!2BJT(^%v2bk4&x+5h!_lD*m}S5GJOWrEu(xGC0cwoOBw2Ks8_1x+jg!{%``y*lzQqp)N& zo+b1UlV?D)5dt1a_{g>J{0J-{{*_~{D|Z>yPD4$Q525JKb}(<3yxX+CW~vu5tZf h`P#pjl6PUk$&VhM+_9`y-P z>$Q^VV041>*G#HY!xpn(#Ia5jdjkh+&_UNMP3IQah#{=o&?s?<%XWvdW!d}3Xv*QS3%Xm>?`O?>wqJVfs%A_3?~U& z|B*|fYGQp8&)G6yV}foS?DJ5@&hH)dBzvGH9#1Cfqa%@f)0d%1Yb;-K!x#$nmkksk z<344@H@AyA?tE(BTgAq52KS$~ypUJe-Et_e&~;W#HI2;HPetzkI5YKCk2 z(pa`WRNgyKQualtz@)f#L!OCt&#E9oc0X2`-RFc7pQ#MX!%_gx66cYc4uK2^eiO=3 zr1cF)ri;Lu=v^$JaO`OIy2j_Xo3hu5DJP||tY+V}&Q*|6b@dR8{LFhK1#ljbD|Bj3 zXf}fwhb=t0)Gw$zq^!rVkhy!7^>{K)Nue`$A#^Senlu-sOtXI>)99`s+P?XM|70=U z=Ysa8lr4qHm%__s$%625G?^K`tA;Wykd2yY2SSJS|A$A#$HMp zfdl)l5U@ma<*Q zoQ7S5%qiJBxa^O3I`+`uQck&s$&umXhCo+=i_&e5Fcp`3JZ&A!gbSj&JmLxI2@ii* z8tB@vhteHRFcsK!(KgGT)Y9SM)H1G14%?EdSp*wY3F|pQ?jI2a19K z&dR;^9xL6-z&LPP)UU2i4`*R#n4@f*!#_dtr&nCuk;|=BmR(f~eMow(jNVKuo!|o1 zZ!tI!zndRRyUoa{RP)OknU>4Pe|3oV;4>Td%uDksQ9j#m-!!y6wTxT$X|;_CUfYKaW*gY>y>OpnAz}*=qY!!k>5v-P_bVZRa+E z7IGS0Bk9F;wyP4(@)9x8h^OJK4?^huIGVlAmm^b*Si5 zP3r;^>c(QzoOJo@yIn^(E%l6db@6>0)prV zQES%Tu5&NUHQo+7$xMlwxCN67TM$i5+On65{m^JJtroK-V=Rj=gor8Mytkj!^X+&3 z&yzgooSdAoVAR-MRvY-=g1igYG(Ba{emMSv^3F`=y6Mrrk2)4oqjuLpC-npXG?0L- z2mrx=FIVs#d-#2;E&gjY2Eh08UzYs|k#71ACmWN|5M!Zyr%FF|rBQwb$h=WVye{=K z8orBdmUj=IQX>(hkg_Hf9hX4Kl1bACKA9{>c15_-|-4ept*WXX$8 zg~H4D?KD9?!6B(|`b(ecDcWRGM>+VK~FRa7&`;P50U;2B)WF$TguYH{h zUCU2>*Co5KxBquGp1Mlf5aW(>=2O&_VXchss7k5H0 zJBRBaV3%!FB^m47?90RoesFUrdkEo} zKvUch1L`BiA=Po7usJ1gY9yp8rOg?EccP()K-iKKxP2m&?GeUlSuOR>gjBse*IG6O z9vuz%oLceYoKgC)^(40LIjl*o=-a}RT?0f_Fm2iIARe@e1_P5y)OW$M>T$>Ty0iGY z=ND~ig*1gHWr)v7ta_4?s>=FexyKI9+4=8FoP+hSA7+@6Jw16XqW1oK;U>LT+Xo=%ygj?w%TCqDr(h^)C5WAn8QXRCcf~lzVqeh`F?+S z{^UvPQ+kcPr`6V95h;9f;&(^yt?c)6N!J_%0Hi+A94L5n{ZO~=&I4CYU%SKC@9H~o z+1}%DItu|{0Sdm60U#NObxNGrDXzEL=2xo;0OGSB7uXQpew? zl%atU4M-jVhDpdlcs;@dn`3>@oB&$kSwH z?Jo?5b_O8&bQtP8mFLyWE9bS$HMq;#p`c7{fTIIG9q%Tc@Fo5dsf83DCk2?{`-=>a z<3A|+!nfi+G1!of?x{a6pkDUO68UMU_^vSLulOTUS^mi8!oJVtnt58TejW@GYzPvT zfzGcUYd>oTMm5XTa&CD|I~JzLLVMl--KOP52NK2X0%7gYSN~ZIXIMcofsc@Zux8*g z)B0e!9@ zEuqm44_!4Fv+ir7H_bAie^iBggVv14oe_KY$TiPT+Sm=#?$D?L@jB^@X92A$HEb=H zNGER*q9z7XNHZZ~e;zOYA36*J94n-+AKF?HU_`kH#K)R85%U6cD;Ff3qa;A(brEvv zG{d`8_09k}{UgqzgRiC9tk87vL$dCc#8X*`+rK2Ie~|e4RzAH=*5%9nf&EgNB&TyY zOWM4X5-;ZXWmRhTTzd;43m)V)s$<^?HO+?;fX1=%$uJ&3%of~SqT6}+#e~kJK8` GlyphRenderer { + let render_ctx = PietGpuRenderContext::new(); + let scale_context = ScaleContext::new(); GlyphRenderer { - render_ctx: PietGpuRenderContext::new(), - scale_context: ScaleContext::new(), + render_ctx, + scale_context, } } - pub unsafe fn add_glyph(&mut self, font_data: &[u8], font_id: u64, glyph_id: u16) { + pub unsafe fn add_glyph( + &mut self, + font_data: &[u8], + font_id: u64, + glyph_id: u16, + transform: [f32; 6], + ) { // This transmute is dodgy because the definition in swash isn't repr(transparent). // I think the best solution is to have a from_u64 method, but we'll work that out // later. let font_id = FontId(std::mem::transmute(font_id)); let encoder = self.make_glyph(font_data, font_id, glyph_id); + const DEFAULT_UPEM: u16 = 2048; + let affine = Affine::new([ + transform[0] as f64, + transform[1] as f64, + transform[2] as f64, + transform[3] as f64, + transform[4] as f64, + transform[5] as f64, + ]) * Affine::scale(1.0 / DEFAULT_UPEM as f64); + self.render_ctx.transform(affine); self.render_ctx.encode_glyph(&encoder); // TODO: don't fill glyph if RGBA self.render_ctx.fill_glyph(0xff_ff_ff_ff); + self.render_ctx.transform(affine.inverse()); + } + + pub fn reset(&mut self) { + self.render_ctx = PietGpuRenderContext::new(); } fn make_glyph(&mut self, font_data: &[u8], font_id: FontId, glyph_id: u16) -> GlyphEncoder { let mut encoder = GlyphEncoder::default(); - let font_ref = FontRef { - data: font_data, - offset: 0, - key: font_id.0, - }; + let font_data = FontDataRef::new(font_data).expect("invalid font"); + let mut font_ref = font_data.get(0).expect("invalid font index"); + font_ref.key = font_id.0; let mut scaler = self.scale_context.builder(font_ref).size(2048.).build(); if let Some(outline) = scaler.scale_outline(glyph_id) { crate::text::append_outline(&mut encoder, outline.verbs(), outline.points()); + } else { + println!("failed to scale"); } encoder } diff --git a/piet-gpu/src/lib.rs b/piet-gpu/src/lib.rs index 47de115..3c1e27f 100644 --- a/piet-gpu/src/lib.rs +++ b/piet-gpu/src/lib.rs @@ -49,6 +49,18 @@ pub fn dump_k1_data(k1_buf: &[u32]) { } } +pub struct RenderConfig { + width: usize, + height: usize, + format: PixelFormat, +} + +// Should we just use the enum from piet-gpu-hal? +pub enum PixelFormat { + A8, + Rgba8, +} + pub struct Renderer { // These sizes are aligned to tile boundaries, though at some point // we'll want to have a good strategy for dealing with odd sizes. @@ -106,15 +118,41 @@ pub struct Renderer { gradients: Image, } +impl RenderConfig { + pub fn new(width: usize, height: usize) -> RenderConfig { + RenderConfig { + width, + height, + format: PixelFormat::Rgba8, + } + } + + pub fn pixel_format(mut self, format: PixelFormat) -> Self { + self.format = format; + self + } +} + impl Renderer { - /// Create a new renderer. pub unsafe fn new( session: &Session, width: usize, height: usize, n_bufs: usize, + ) -> Result { + let config = RenderConfig::new(width, height); + Self::new_from_config(session, config, n_bufs) + } + + /// Create a new renderer. + pub unsafe fn new_from_config( + session: &Session, + config: RenderConfig, + n_bufs: usize, ) -> Result { // For now, round up to tile alignment + let width = config.width; + let height = config.height; let width = width + (width.wrapping_neg() & (TILE_W - 1)); let height = height + (height.wrapping_neg() & (TILE_W - 1)); let dev = BufferUsage::STORAGE | BufferUsage::COPY_DST; @@ -126,7 +164,11 @@ impl Renderer { .map(|_| session.create_buffer(8 * 1024 * 1024, host_upload).unwrap()) .collect::>(); - let image_dev = session.create_image2d(width as u32, height as u32)?; + let image_format = match config.format { + PixelFormat::A8 => piet_gpu_hal::ImageFormat::A8, + PixelFormat::Rgba8 => piet_gpu_hal::ImageFormat::Rgba8, + }; + let image_dev = session.create_image2d(width as u32, height as u32, image_format)?; // Note: this must be updated when the config struct size changes. const CONFIG_BUFFER_SIZE: u64 = std::mem::size_of::() as u64; @@ -211,7 +253,10 @@ impl Renderer { .collect(); let gradients = Self::make_gradient_image(&session); - let k4_code = include_shader!(session, "../shader/gen/kernel4"); + let k4_code = match config.format { + PixelFormat::A8 => include_shader!(session, "../shader/gen/kernel4_gray"), + PixelFormat::Rgba8 => include_shader!(session, "../shader/gen/kernel4"), + }; let k4_pipeline = session.create_compute_pipeline( k4_code, &[ @@ -428,7 +473,8 @@ impl Renderer { return Err("unsupported image format".into()); } let buffer = session.create_buffer_init(&buf, BufferUsage::COPY_SRC)?; - let image = session.create_image2d(width.try_into()?, height.try_into()?)?; + const RGBA: piet_gpu_hal::ImageFormat = piet_gpu_hal::ImageFormat::Rgba8; + let image = session.create_image2d(width.try_into()?, height.try_into()?, RGBA)?; let mut cmd_buf = session.cmd_buf()?; cmd_buf.begin(); cmd_buf.image_barrier(&image, ImageLayout::Undefined, ImageLayout::BlitDst); @@ -464,8 +510,13 @@ impl Renderer { fn make_gradient_image(session: &Session) -> Image { unsafe { + const RGBA: piet_gpu_hal::ImageFormat = piet_gpu_hal::ImageFormat::Rgba8; session - .create_image2d(gradient::N_SAMPLES as u32, gradient::N_GRADIENTS as u32) + .create_image2d( + gradient::N_SAMPLES as u32, + gradient::N_GRADIENTS as u32, + RGBA, + ) .unwrap() } }