From d948126c1662ad71448e8cca14c6df21328c48cb Mon Sep 17 00:00:00 2001
From: Raph Levien <raph.levien@gmail.com>
Date: Wed, 8 Dec 2021 10:42:35 -0800
Subject: [PATCH] Adjust workgroup sizes

Make max workgroup size 256 and respect LG_WG_FACTOR.

Because the monoid scans only support a height of 2, this will reduce
the maximum scene complexity we can render. But it also increases
compatibility. Supporting larger scans is a TODO.
---
 piet-gpu/shader/binning.comp              |   2 +-
 piet-gpu/shader/build.ninja               |   6 +--
 piet-gpu/shader/draw_leaf.comp            |   2 +-
 piet-gpu/shader/draw_reduce.comp          |   2 +-
 piet-gpu/shader/draw_scan.comp            |   3 +-
 piet-gpu/shader/gen/binning.hlsl          |  38 +++++++--------
 piet-gpu/shader/gen/binning.msl           |  16 +++----
 piet-gpu/shader/gen/binning.spv           | Bin 16416 -> 16400 bytes
 piet-gpu/shader/gen/draw_leaf.dxil        | Bin 6880 -> 6880 bytes
 piet-gpu/shader/gen/draw_leaf.hlsl        |  54 +++++++++++-----------
 piet-gpu/shader/gen/draw_leaf.msl         |  32 ++++++-------
 piet-gpu/shader/gen/draw_leaf.spv         | Bin 38560 -> 38560 bytes
 piet-gpu/shader/gen/draw_reduce.dxil      | Bin 3940 -> 3940 bytes
 piet-gpu/shader/gen/draw_reduce.hlsl      |  20 ++++----
 piet-gpu/shader/gen/draw_reduce.msl       |  14 +++---
 piet-gpu/shader/gen/draw_reduce.spv       | Bin 6872 -> 6856 bytes
 piet-gpu/shader/gen/draw_root.dxil        | Bin 3944 -> 3944 bytes
 piet-gpu/shader/gen/draw_root.hlsl        |  14 +++---
 piet-gpu/shader/gen/draw_root.msl         |  12 ++---
 piet-gpu/shader/gen/draw_root.spv         | Bin 4684 -> 4668 bytes
 piet-gpu/shader/gen/pathseg.dxil          | Bin 9596 -> 9592 bytes
 piet-gpu/shader/gen/pathseg.hlsl          |  48 +++++++++----------
 piet-gpu/shader/gen/pathseg.msl           |  38 +++++++--------
 piet-gpu/shader/gen/pathseg.spv           | Bin 34748 -> 34732 bytes
 piet-gpu/shader/gen/pathtag_reduce.dxil   | Bin 4700 -> 4644 bytes
 piet-gpu/shader/gen/pathtag_reduce.hlsl   |  18 ++++----
 piet-gpu/shader/gen/pathtag_reduce.msl    |  16 +++----
 piet-gpu/shader/gen/pathtag_reduce.spv    | Bin 7836 -> 7820 bytes
 piet-gpu/shader/gen/pathtag_root.dxil     | Bin 4716 -> 4716 bytes
 piet-gpu/shader/gen/pathtag_root.hlsl     |  20 ++++----
 piet-gpu/shader/gen/pathtag_root.msl      |  18 ++++----
 piet-gpu/shader/gen/pathtag_root.spv      | Bin 5852 -> 5836 bytes
 piet-gpu/shader/gen/tile_alloc.hlsl       |  32 ++++++-------
 piet-gpu/shader/gen/tile_alloc.msl        |  22 ++++-----
 piet-gpu/shader/gen/tile_alloc.spv        | Bin 15176 -> 15160 bytes
 piet-gpu/shader/gen/transform_leaf.dxil   | Bin 5664 -> 5664 bytes
 piet-gpu/shader/gen/transform_leaf.hlsl   |  34 +++++++-------
 piet-gpu/shader/gen/transform_leaf.msl    |  14 +++---
 piet-gpu/shader/gen/transform_leaf.spv    | Bin 12524 -> 12508 bytes
 piet-gpu/shader/gen/transform_reduce.dxil | Bin 4696 -> 4696 bytes
 piet-gpu/shader/gen/transform_reduce.hlsl |  18 ++++----
 piet-gpu/shader/gen/transform_reduce.msl  |  14 +++---
 piet-gpu/shader/gen/transform_reduce.spv  | Bin 7876 -> 7860 bytes
 piet-gpu/shader/gen/transform_root.dxil   | Bin 4824 -> 4824 bytes
 piet-gpu/shader/gen/transform_root.hlsl   |  14 +++---
 piet-gpu/shader/gen/transform_root.msl    |  12 ++---
 piet-gpu/shader/gen/transform_root.spv    | Bin 5280 -> 5336 bytes
 piet-gpu/shader/pathseg.comp              |   2 +-
 piet-gpu/shader/pathtag_reduce.comp       |   4 +-
 piet-gpu/shader/pathtag_scan.comp         |   3 +-
 piet-gpu/shader/setup.h                   |   2 +
 piet-gpu/shader/tile_alloc.comp           |   2 +-
 piet-gpu/shader/transform_leaf.comp       |   2 +-
 piet-gpu/shader/transform_reduce.comp     |   2 +-
 piet-gpu/shader/transform_scan.comp       |   5 +-
 piet-gpu/src/encoder.rs                   |  22 ++++-----
 piet-gpu/src/render_ctx.rs                |  21 +--------
 piet-gpu/src/stages.rs                    |   8 ++--
 piet-gpu/src/stages/draw.rs               |   4 +-
 piet-gpu/src/stages/path.rs               |  11 +++--
 piet-gpu/src/stages/transform.rs          |   4 +-
 piet-gpu/src/text.rs                      |   3 +-
 tests/src/draw.rs                         |   3 +-
 tests/src/path.rs                         |   3 +-
 tests/src/transform.rs                    |   3 +-
 65 files changed, 313 insertions(+), 324 deletions(-)

diff --git a/piet-gpu/shader/binning.comp b/piet-gpu/shader/binning.comp
index 313310e..c2b81fd 100644
--- a/piet-gpu/shader/binning.comp
+++ b/piet-gpu/shader/binning.comp
@@ -84,7 +84,7 @@ void main() {
     if (x0 == x1) y1 = y0;
     int x = x0, y = y0;
     uint my_slice = gl_LocalInvocationID.x / 32;
-    uint my_mask = 1 << (gl_LocalInvocationID.x & 31);
+    uint my_mask = 1u << (gl_LocalInvocationID.x & 31);
     while (y < y1) {
         atomicOr(bitmaps[my_slice][y * width_in_bins + x], my_mask);
         x++;
diff --git a/piet-gpu/shader/build.ninja b/piet-gpu/shader/build.ninja
index 6ed2140..6a59f59 100644
--- a/piet-gpu/shader/build.ninja
+++ b/piet-gpu/shader/build.ninja
@@ -65,7 +65,7 @@ build gen/transform_reduce.hlsl: hlsl gen/transform_reduce.spv
 build gen/transform_reduce.dxil: dxil gen/transform_reduce.hlsl
 build gen/transform_reduce.msl: msl gen/transform_reduce.spv
 
-build gen/transform_root.spv: glsl transform_scan.comp
+build gen/transform_root.spv: glsl transform_scan.comp | setup.h
   flags = -DROOT
 build gen/transform_root.hlsl: hlsl gen/transform_root.spv
 build gen/transform_root.dxil: dxil gen/transform_root.hlsl
@@ -81,7 +81,7 @@ build gen/pathtag_reduce.hlsl: hlsl gen/pathtag_reduce.spv
 build gen/pathtag_reduce.dxil: dxil gen/pathtag_reduce.hlsl
 build gen/pathtag_reduce.msl: msl gen/pathtag_reduce.spv
 
-build gen/pathtag_root.spv: glsl pathtag_scan.comp | pathtag.h
+build gen/pathtag_root.spv: glsl pathtag_scan.comp | pathtag.h setup.h
   flags = -DROOT
 build gen/pathtag_root.hlsl: hlsl gen/pathtag_root.spv
 build gen/pathtag_root.dxil: dxil gen/pathtag_root.hlsl
@@ -102,7 +102,7 @@ build gen/draw_reduce.hlsl: hlsl gen/draw_reduce.spv
 build gen/draw_reduce.dxil: dxil gen/draw_reduce.hlsl
 build gen/draw_reduce.msl: msl gen/draw_reduce.spv
 
-build gen/draw_root.spv: glsl draw_scan.comp | drawtag.h
+build gen/draw_root.spv: glsl draw_scan.comp | drawtag.h setup.h
   flags = -DROOT
 build gen/draw_root.hlsl: hlsl gen/draw_root.spv
 build gen/draw_root.dxil: dxil gen/draw_root.hlsl
diff --git a/piet-gpu/shader/draw_leaf.comp b/piet-gpu/shader/draw_leaf.comp
index 85d9528..5de2652 100644
--- a/piet-gpu/shader/draw_leaf.comp
+++ b/piet-gpu/shader/draw_leaf.comp
@@ -11,7 +11,7 @@
 #include "setup.h"
 
 #define N_ROWS 8
-#define LG_WG_SIZE 9
+#define LG_WG_SIZE (7 + LG_WG_FACTOR)
 #define WG_SIZE (1 << LG_WG_SIZE)
 #define PARTITION_SIZE (WG_SIZE * N_ROWS)
 
diff --git a/piet-gpu/shader/draw_reduce.comp b/piet-gpu/shader/draw_reduce.comp
index fe9ab2c..68d43e9 100644
--- a/piet-gpu/shader/draw_reduce.comp
+++ b/piet-gpu/shader/draw_reduce.comp
@@ -9,7 +9,7 @@
 #include "setup.h"
 
 #define N_ROWS 8
-#define LG_WG_SIZE 9
+#define LG_WG_SIZE (7 + LG_WG_FACTOR)
 #define WG_SIZE (1 << LG_WG_SIZE)
 #define PARTITION_SIZE (WG_SIZE * N_ROWS)
 
diff --git a/piet-gpu/shader/draw_scan.comp b/piet-gpu/shader/draw_scan.comp
index d883671..2afc9ba 100644
--- a/piet-gpu/shader/draw_scan.comp
+++ b/piet-gpu/shader/draw_scan.comp
@@ -5,10 +5,11 @@
 #version 450
 #extension GL_GOOGLE_include_directive : enable
 
+#include "setup.h"
 #include "drawtag.h"
 
 #define N_ROWS 8
-#define LG_WG_SIZE 9
+#define LG_WG_SIZE (7 + LG_WG_FACTOR)
 #define WG_SIZE (1 << LG_WG_SIZE)
 #define PARTITION_SIZE (WG_SIZE * N_ROWS)
 
diff --git a/piet-gpu/shader/gen/binning.hlsl b/piet-gpu/shader/gen/binning.hlsl
index 2b0901e..b13db37 100644
--- a/piet-gpu/shader/gen/binning.hlsl
+++ b/piet-gpu/shader/gen/binning.hlsl
@@ -248,11 +248,11 @@ void comp_main()
     int x = x0;
     int y = y0;
     uint my_slice = gl_LocalInvocationID.x / 32u;
-    uint my_mask = uint(1 << int(gl_LocalInvocationID.x & 31u));
+    uint my_mask = 1u << (gl_LocalInvocationID.x & 31u);
     while (y < y1)
     {
-        uint _438;
-        InterlockedOr(bitmaps[my_slice][(uint(y) * width_in_bins) + uint(x)], my_mask, _438);
+        uint _437;
+        InterlockedOr(bitmaps[my_slice][(uint(y) * width_in_bins) + uint(x)], my_mask, _437);
         x++;
         if (x == x1)
         {
@@ -274,8 +274,8 @@ void comp_main()
     if (element_count != 0u)
     {
         uint param_7 = element_count * 4u;
-        MallocResult _488 = malloc(param_7);
-        MallocResult chunk = _488;
+        MallocResult _487 = malloc(param_7);
+        MallocResult chunk = _487;
         chunk_alloc = chunk.alloc;
         sh_chunk_alloc[gl_LocalInvocationID.x] = chunk_alloc;
         if (chunk.failed)
@@ -284,31 +284,31 @@ void comp_main()
         }
     }
     uint out_ix = (_253.Load(20) >> uint(2)) + (((my_partition * 256u) + gl_LocalInvocationID.x) * 2u);
-    Alloc _517;
-    _517.offset = _253.Load(20);
+    Alloc _516;
+    _516.offset = _253.Load(20);
     Alloc param_8;
-    param_8.offset = _517.offset;
+    param_8.offset = _516.offset;
     uint param_9 = out_ix;
     uint param_10 = element_count;
     write_mem(param_8, param_9, param_10);
-    Alloc _529;
-    _529.offset = _253.Load(20);
+    Alloc _528;
+    _528.offset = _253.Load(20);
     Alloc param_11;
-    param_11.offset = _529.offset;
+    param_11.offset = _528.offset;
     uint param_12 = out_ix + 1u;
     uint param_13 = chunk_alloc.offset;
     write_mem(param_11, param_12, param_13);
     GroupMemoryBarrierWithGroupSync();
-    bool _544;
+    bool _543;
     if (!sh_alloc_failed)
     {
-        _544 = _84.Load(4) != 0u;
+        _543 = _84.Load(4) != 0u;
     }
     else
     {
-        _544 = sh_alloc_failed;
+        _543 = sh_alloc_failed;
     }
-    if (_544)
+    if (_543)
     {
         return;
     }
@@ -327,11 +327,11 @@ void comp_main()
             }
             Alloc out_alloc = sh_chunk_alloc[bin_ix];
             uint out_offset = out_alloc.offset + (idx * 4u);
-            BinInstanceRef _606 = { out_offset };
-            BinInstance _608 = { element_ix };
+            BinInstanceRef _605 = { out_offset };
+            BinInstance _607 = { element_ix };
             Alloc param_14 = out_alloc;
-            BinInstanceRef param_15 = _606;
-            BinInstance param_16 = _608;
+            BinInstanceRef param_15 = _605;
+            BinInstance param_16 = _607;
             BinInstance_write(param_14, param_15, param_16);
         }
         x++;
diff --git a/piet-gpu/shader/gen/binning.msl b/piet-gpu/shader/gen/binning.msl
index f6e0505..42a11ee 100644
--- a/piet-gpu/shader/gen/binning.msl
+++ b/piet-gpu/shader/gen/binning.msl
@@ -260,10 +260,10 @@ kernel void main0(constant uint* spvBufferSizeConstants [[buffer(25)]], device M
     int x = x0;
     int y = y0;
     uint my_slice = gl_LocalInvocationID.x / 32u;
-    uint my_mask = uint(1 << int(gl_LocalInvocationID.x & 31u));
+    uint my_mask = 1u << (gl_LocalInvocationID.x & 31u);
     while (y < y1)
     {
-        uint _438 = atomic_fetch_or_explicit((threadgroup atomic_uint*)&bitmaps[my_slice][(uint(y) * width_in_bins) + uint(x)], my_mask, memory_order_relaxed);
+        uint _437 = atomic_fetch_or_explicit((threadgroup atomic_uint*)&bitmaps[my_slice][(uint(y) * width_in_bins) + uint(x)], my_mask, memory_order_relaxed);
         x++;
         if (x == x1)
         {
@@ -285,8 +285,8 @@ kernel void main0(constant uint* spvBufferSizeConstants [[buffer(25)]], device M
     if (element_count != 0u)
     {
         uint param_7 = element_count * 4u;
-        MallocResult _488 = malloc(param_7, v_84, v_84BufferSize);
-        MallocResult chunk = _488;
+        MallocResult _487 = malloc(param_7, v_84, v_84BufferSize);
+        MallocResult chunk = _487;
         chunk_alloc = chunk.alloc;
         sh_chunk_alloc[gl_LocalInvocationID.x] = chunk_alloc;
         if (chunk.failed)
@@ -306,16 +306,16 @@ kernel void main0(constant uint* spvBufferSizeConstants [[buffer(25)]], device M
     uint param_13 = chunk_alloc.offset;
     write_mem(param_11, param_12, param_13, v_84, v_84BufferSize);
     threadgroup_barrier(mem_flags::mem_threadgroup);
-    bool _544;
+    bool _543;
     if (!bool(sh_alloc_failed))
     {
-        _544 = v_84.mem_error != 0u;
+        _543 = v_84.mem_error != 0u;
     }
     else
     {
-        _544 = bool(sh_alloc_failed);
+        _543 = bool(sh_alloc_failed);
     }
-    if (_544)
+    if (_543)
     {
         return;
     }
diff --git a/piet-gpu/shader/gen/binning.spv b/piet-gpu/shader/gen/binning.spv
index 5ec7aec1bbb810f495c2e8ea6b4b0f14dd863c98..17043bc0bc5b911798d9a4c360ac6de0b53c913c 100644
GIT binary patch
delta 4554
zcmZvfXK+?U6o&Upp#&nt1W=;_B3Qs)Q9)v%S+Qcl7%)f#5=UXiA3B#&|M<fhXFkT!
zah$RD79%S5j=lG8?226}D(drmxhH(XcxN`}ecyBT?C#mUd%s+@q;ScSLTzQmsJ<x`
zQcdcY&g`tcV!+Xl6;f(l(7Ld-eNpqAj>YYZQcBgSGChvoR>D>3b?|y$b$VYuZ*lwl
z=9UEuI_3~xl|Dc_v*oOog((%kA2534=Mj_V_-ehM_&#Z6v1#D28LLn`7B6aUJI7J0
zQbk2?mJ;u~0q?&7-)I9~U%Y(a@U|ULXScQIS*z0?a7S6#a`t?tL*~jf39aq4bHLT<
zsNCPP8vWJjxLoU)JNN8XisI}}Ma$P%;<Jk*Ya6=`t({j<=;&HDsJVKAm8xQd&H4Ks
zxCT5JI|BZF?o;14=PPqf|Ifu4L+a~Rqg2yt95pI{JGvGRxu8&ZtoZHF#=4<6EIxZ!
zO3}ACWZ3YmKkTD2^(46E?~Q_6{&H^NgPYaD2iIEoa&F;+8*kxr-lqiWi_4jcm5(s3
zl`rR3J-BJCd~mH*53b*e2j3q)yNox%=j6PJY8{(W*Xf&d)?|HfT*LZiVz0(dWV^4#
zs;x~G>00cX+^(YV*TZk1`RKnf_iHEMpN4)D_|h_dGk6-Ulg;Xf{}!+@=vn%cP72r2
zOzvP}lR3Ft!I$OS<hOxMes1o+9jq4dcYxbz7c!Xr?*tp89uvF^yp(os%6bOrgeWe_
zJKYVoqq^nV6x&5>pgAG^G2lI5EA|V;cfe&}W3A{3H2pE~y<i7+92aw622q>^vGh9c
zqeX}N!D_*mgQwDN$xHbF*ckQI^nVEJV^#G%NK><_;>hzbcrvXm22}{>BVbD$vwRe+
z7L(Mc=Ke==t=zNWVA}HYbA9pR(Zez?P;26-UYcT~p2&M?J!MZ}pQicfe<t^9ouzeN
z0XD~7c`2R+t8GMc`SK@e=5j{h&3ZdOhiK;}I#*$zFJ8S>efEN=r1ot2i{J*D(fVTO
zm%#d7Vi>paWic(z<SSsch<_FA(&Ko&239l9Qs_flKfN~)%BS~cX?ky=zDbL;Z-ea`
zC-fb#m#V#uHMN>?R?&)D32o$f7wkwewfDekvDf<4=%4j&lO;8|-?7-v56kdJ;Pa@<
zcwPx_^^f6$u*Msw-~L|D`eHJlfL+k7`IdePR&$KiDXUD~5Fhtf=VvrES0#=_Uw};%
zm%x``wQl-4@K<29@eEXl_BFTxtFGTKf||#}tMwc3=i~S@9pWIa>RY&uScyE&(+{N$
z`w`^_TJ-n{?3#`G8T$)uCCxaq$c@q8pAmiqI~6y>9A2Kk!PR}+|G@fK4)xz@YL)|>
zm6DD^a5Fps;nd2z^I&Zgo5g1#tpTsYx>P@~v1H%5wZ_w)V1n{L!D<(QqxOHn6Yav(
z*c0?OL~$@J>a(s?pMOx+(c<v^3pQFkwvJgQ^-i)zTmj9Y^;x%7U^QQa{lJyr@+CO6
zwS-rp>zmOt{VY{O6n&hUpU*yUzwoZx<Z|O;ef_`=sLwCIJbX1^&z(Mh1;~8~FlhjP
zMqo?79>D#1p?W(GL@;TruokZ7M<iA?2>ySo8jP+lUsb9DtNFyDHir8%Bo-yFo5>$%
zT90+W2p$6Vn3nJSP&9Sp8?f?N*)Z@326F&&H58YRZOk?SdG>Ci_#WI%;c6b-My!tq
zSA8>@ng>_xOgxL5gO9<^qD76jKvR!PVoR`Ed`t2?St;Qk6yL#z6X4+8Ovnxn;zUMb
zeVm9fBWP+)M7%B)N?6Sac-xNxPbQD0Sc;X0Z)@<Bl8@W!Qh<VQ+GwmvwSM8Y!EQ_Q
z(Z5~p*G97K!R81)7CsFgzoW*%jZt?2@-Z}XugOR%iOxHMO%~s%o#1NveU*0xx8slZ
zZWp*Q>XBntuq%!nyTR4s)b0+ppL%9_{E4d=gX{^uj0C~=f*TWGu)V<z_>I#SapS@I
z%<0c$dHD7Ls|DW|e5IG)iu{t=55XWEacCxh)qB1ym1%#lk7d+1k*4OBi6eOvxO}h!
z;0%_|&&N6tt~QhLdna%Zf*p0nPlkiR>QTzcU^Q1b2`i65r+}^6Z4BaJo(fls$xjFS
zIC*{3Xwx(PID1JXIs|N@nC+o(HE)@SISgLDC5OZPQDLmUh&uwTFZhvQCmv@Zn*mWY
zFjl4JX?Mm>P8$i12KNly^Km=|?#EHxm>MQ|ELh!qM`7i0YaI_(%Wti$bgE|}=!ij1
z06WV#<tKvGx=9{0KM8I>eGYUSRxNUz3|2dmNlm2REU;SmPXUi9u>S^3BA^*TF;>w6
zcBUpj6)P9d#+pn#2W!8$V^0H5CYSNX>9>+`NVVqof@%dX%_yBKQkaXN=rw9`-(j`r
zaJn7Pg3p6bg~xvxZE$1My*1}!eLQgb&Y-EerDF4B4!i(TepSk2kcD6`q?lwoTrDQ4
gPt8?Ea;<z@tb=LgqjSCu)+>AOROp(r>wC5T0T%onmjD0&

delta 4567
zcmZvfXK<B86o&UExe)>}L;@(NU;!Hz?5GIWf?&rAD#i#wBaooLj6ZZfM*ZVIXD;J7
z<2Yki>==9Rz4u-ayQru%74>;;zJpiBZ)S7e_dRFN?w;Lmzne=I6s}xQXsoT-YJ-#t
zsUh`BXD@ENyw8D;7gB1U)jqqub6)GTuKAtwQcA^Co1Q@LsNlNvCiq`pF?}dMbAIQ{
z*0x!*x~7%?=rerOCunop&S{&S_gNY-Isf`U{`$16T<SZhv<kIr{=C+Xa~-uV)z<ts
zORx2K!+N}_yr}QcjsfebJHb!y=*+Vg)1GigncX&bCex{`uY}ey{akP{9h>_rtEu>l
z>BL;?nlWQ;J4JDJZRJsoLyDs5-KEC4HHA6dcQ>^b$5@{_`fbYJ@4*e=e%N8~A9A1i
zzByl(Yx<X$Ck^P|yc(rQuaT)f6~J@4=MT8BQ0_ITZ!`N<*$k}-29}2p8k)&ICNR|A
zaJL}%$#AzIxYjMG<`zHrK>QXzxYpuVbBmwz`ecWu^7Bl^>Q{$Z{c3LYtGSgAZWgN_
zTx;cn>u;;#W8u@PcqyeCHe<@hk2Yiz@m<Z5red$co=36j!PkJ*_vch!ft9bJ+qKx$
zxm`)@_3#^LKKgIW{n|13&A$kIX%)W-Jdw6I6*FF`lx~I?q~muw$*-fC+`+^qb7HrE
zFUz^fZv~q?CVm@OE#hwnchWB6sM!AwurcamQr0uW#gK)xNOC7wZ9(3t8*E4Q5@|NU
z-bHJnIU)Tq;N4&=ww(QNz<a>PP6kh*>5t0Xn?W4dfn3df;CpFqho#r279H*fs|9}m
zJb`u-{jKDKU}Mx*(f?trk5$$85KYahinB=c2xL4h29<~JQLrVBNj?Tvi%IHJbN?f`
zR_@twFm3huace(Ib>$0(H)r0aR>rZtLB+;C3AV4+v-ULh8Jds&XLG;SxmxY#z~;Ck
zFUa#?wGC+&LH-oYT;>K3Uf1~rL_3%0T!($JeASlyvzJ5_wP)jB2Di|R))zy+0@n93
z!?>5Pg4N<sy#`i`_}9TMJr3C$U^U|`g+8=@r}q{@_4MAZOz$1kw`q~~U9er_oW2M4
zuC=$ZrdBh~Dq2x1p^Y5xgB>ZR_5oNec3htt{j+u3WJyi_5jM8-<0|Y;dp>m;&5Pk>
z|0%o)YrJv#?eB+DUrgpRunW35-_p;)YL2lgWwq%Gh>!cLa|un&Rf!|fmtYgcRqz#9
zZ3%rH_-nA*Xa;IV`v%;CRoA~1tLE|Wj{O$AbR>Tcg*b?-`VOumRw9q{^kZeienR<?
z7Cn9jyJn-7V}GG7qZwxwxiR{CGs3T6r{ZRq!}tC-xVrEBO019NP+vh)vmD^8l5{Hs
zH^UPUPOUs@r!BEr#~TM;i*>21!LelDx!>(adx8nd{{X994366W2~V_((vY5@wGhRA
zw5ZQtmHPaR@)s=*A2UiutH;(A(9{=WBCZDR(E6;KUurd9g#EyE;OZqfwKeoFqU)Q~
zGyN>pLlk|Sn%~t8;C|{|x5?$k#rk@K9Z;V?0p#K91NPkM^XGuvmjILc@@E*<-1ckW
zg<99K5y7Of!VTeSeoA6h{owz%s%CV3`Kr=JU^SmuRDXCAEfyu;Xfl7CX$#f?BX}U#
zV_LoQgV5BC_ZN*kR<<#C7=t;0xdxOM?Jy+U1mxMfk>Y!B2gB7oxSL{qJh<vZXlfo@
zu`}^3ZU#OMJCznS-W*LmE{UOFwfL6gd9q5v6%^mWhY{f59&|f6h!fco>*GX>*@C9#
zM8s=3$znAp;BCJ(cszM5#X_t+e8a(qRD9f4mjV=g(?(!Tr1cB8Ep|JakN)j*zc!NX
z05(VPk?@Ideoa*twIkdZbr&GthGy>78A&D4d1tW6;v2OKTur~P@+fd8{&??pg&U(D
zId%iP;>fW(TrE!R9$@>aXO_p`x{5K#Uf|0}5PWaAG4TZ(4Q|12oW6+L2dvMW{$7@c
zZ(p!l@cqD7c=<is^OmPE2nOkhL$g0vz30nPo5q2CETg`$G&Q$O9LWy=R}Xd|oWars
z`B(?R)h09kx(OVNU`L(tlVLnqJxX~9Sj`odu<{sm0@$kElJEOOxLQp9FmOy>-=Va_
zGX6MwNhCTPY@(R$5pXqcnTR<OUcDtp!TnKTtiFgl8muq)B(M{YvydGFQ8X}CrRHgO
z#!gNf362Bz4BYc^JRa`HQQepZCV2u_-F(Ml<#B6yo2liu7MD96rXc8uK~4fY%Q)pH
zgVmN08#6xzZa;kvbRt$Qa!duQ9qn16->G1=@VA0T6xe?Q#u0EDf?}-Vbg(luc^g(P
zo`yA<cska8amTiU$CJx=<Mdm}aab#<HNO|s4DiB?(zzmqGY}NLMosQJtQH;4v;$i3
z4)_Fk{I78q+!%Fl%~@C<-xhr{X=+YHY@W*ho3kO+SEW1#=>&Tr#U#7nYB5QDYOXSp
bYvtQv9ZV}9k@Ib_UfFx6Lib*~zTfx{9KRi3

diff --git a/piet-gpu/shader/gen/draw_leaf.dxil b/piet-gpu/shader/gen/draw_leaf.dxil
index 86b37e940da57e098366edaadf2139a3bc786b49..17bfd04821a365b1d94408450ff3b37ff214016d 100644
GIT binary patch
delta 125
zcmV-@0D}MEHQ+TAL|8&YR7Ec4m9uUObm<;+6)hUxkrasm0kNKS0Ro@^vlRl@0SAzv
z00RmTq?7XlloS}62?7vpNCxF0prIiE7GE(iGdQ!G1mOb^PRHZaEQZ301N{Zz-U>rs
flZ*)W0-yl1g$d&U0+4{SY73qP0)Q5?3K7-`0z)kk

delta 125
zcmV-@0D}MEHQ+TAL|8&Y4eJ2Z1yB{%;uDeu`44KTkrasm0<oTT0RkWavlRl@0S6GF
z00RmTsFU*ploXhn2?7vpNCxF0prIiE7GE(iGdZ)H1mOb^^osC8MrN8Y6m1qh9<=`H
flZ*)W0w4jig$d&U0uX_-Y73qP0>Bls3K7-`uM;b9

diff --git a/piet-gpu/shader/gen/draw_leaf.hlsl b/piet-gpu/shader/gen/draw_leaf.hlsl
index e3cb387..d0bef52 100644
--- a/piet-gpu/shader/gen/draw_leaf.hlsl
+++ b/piet-gpu/shader/gen/draw_leaf.hlsl
@@ -151,7 +151,7 @@ struct Config
     uint pathseg_offset;
 };
 
-static const uint3 gl_WorkGroupSize = uint3(512u, 1u, 1u);
+static const uint3 gl_WorkGroupSize = uint3(256u, 1u, 1u);
 
 static const DrawMonoid _418 = { 0u, 0u };
 static const DrawMonoid _442 = { 1u, 0u };
@@ -159,8 +159,8 @@ static const DrawMonoid _444 = { 1u, 1u };
 
 RWByteAddressBuffer _201 : register(u0, space0);
 ByteAddressBuffer _225 : register(t2, space0);
-ByteAddressBuffer _1005 : register(t3, space0);
-ByteAddressBuffer _1039 : register(t1, space0);
+ByteAddressBuffer _1004 : register(t3, space0);
+ByteAddressBuffer _1038 : register(t1, space0);
 
 static uint3 gl_WorkGroupID;
 static uint3 gl_LocalInvocationID;
@@ -172,7 +172,7 @@ struct SPIRV_Cross_Input
     uint3 gl_GlobalInvocationID : SV_DispatchThreadID;
 };
 
-groupshared DrawMonoid sh_scratch[512];
+groupshared DrawMonoid sh_scratch[256];
 
 ElementTag Element_tag(ElementRef ref)
 {
@@ -558,7 +558,7 @@ void comp_main()
         local[i] = agg;
     }
     sh_scratch[gl_LocalInvocationID.x] = agg;
-    for (uint i_1 = 0u; i_1 < 9u; i_1++)
+    for (uint i_1 = 0u; i_1 < 8u; i_1++)
     {
         GroupMemoryBarrierWithGroupSync();
         if (gl_LocalInvocationID.x >= (1u << i_1))
@@ -575,11 +575,11 @@ void comp_main()
     DrawMonoid row = tag_monoid_identity();
     if (gl_WorkGroupID.x > 0u)
     {
-        DrawMonoid _1011;
-        _1011.path_ix = _1005.Load((gl_WorkGroupID.x - 1u) * 8 + 0);
-        _1011.clip_ix = _1005.Load((gl_WorkGroupID.x - 1u) * 8 + 4);
-        row.path_ix = _1011.path_ix;
-        row.clip_ix = _1011.clip_ix;
+        DrawMonoid _1010;
+        _1010.path_ix = _1004.Load((gl_WorkGroupID.x - 1u) * 8 + 0);
+        _1010.clip_ix = _1004.Load((gl_WorkGroupID.x - 1u) * 8 + 4);
+        row.path_ix = _1010.path_ix;
+        row.clip_ix = _1010.clip_ix;
     }
     if (gl_LocalInvocationID.x > 0u)
     {
@@ -588,9 +588,9 @@ void comp_main()
         row = combine_tag_monoid(param_10, param_11);
     }
     uint out_ix = gl_GlobalInvocationID.x * 8u;
-    uint out_base = (_1039.Load(44) >> uint(2)) + (out_ix * 2u);
-    AnnotatedRef _1055 = { _1039.Load(32) + (out_ix * 40u) };
-    AnnotatedRef out_ref = _1055;
+    uint out_base = (_1038.Load(44) >> uint(2)) + (out_ix * 2u);
+    AnnotatedRef _1054 = { _1038.Load(32) + (out_ix * 40u) };
+    AnnotatedRef out_ref = _1054;
     float4 mat;
     float2 translate;
     AnnoColor anno_fill;
@@ -617,7 +617,7 @@ void comp_main()
         tag_word = Element_tag(param_16).tag;
         if (((tag_word == 4u) || (tag_word == 5u)) || (tag_word == 6u))
         {
-            uint bbox_offset = (_1039.Load(40) >> uint(2)) + (6u * (m.path_ix - 1u));
+            uint bbox_offset = (_1038.Load(40) >> uint(2)) + (6u * (m.path_ix - 1u));
             float bbox_l = float(_201.Load(bbox_offset * 4 + 8)) - 32768.0f;
             float bbox_t = float(_201.Load((bbox_offset + 1u) * 4 + 8)) - 32768.0f;
             float bbox_r = float(_201.Load((bbox_offset + 2u) * 4 + 8)) - 32768.0f;
@@ -628,7 +628,7 @@ void comp_main()
             if ((linewidth >= 0.0f) || (tag_word == 5u))
             {
                 uint trans_ix = _201.Load((bbox_offset + 5u) * 4 + 8);
-                uint t = (_1039.Load(36) >> uint(2)) + (6u * trans_ix);
+                uint t = (_1038.Load(36) >> uint(2)) + (6u * trans_ix);
                 mat = asfloat(uint4(_201.Load(t * 4 + 8), _201.Load((t + 1u) * 4 + 8), _201.Load((t + 2u) * 4 + 8), _201.Load((t + 3u) * 4 + 8)));
                 if (tag_word == 5u)
                 {
@@ -649,9 +649,9 @@ void comp_main()
                     anno_fill.bbox = bbox;
                     anno_fill.linewidth = linewidth;
                     anno_fill.rgba_color = fill.rgba_color;
-                    Alloc _1258;
-                    _1258.offset = _1039.Load(32);
-                    param_18.offset = _1258.offset;
+                    Alloc _1257;
+                    _1257.offset = _1038.Load(32);
+                    param_18.offset = _1257.offset;
                     AnnotatedRef param_19 = out_ref;
                     uint param_20 = fill_mode;
                     AnnoColor param_21 = anno_fill;
@@ -674,9 +674,9 @@ void comp_main()
                     anno_lin.line_x = line_x;
                     anno_lin.line_y = line_y;
                     anno_lin.line_c = -((p0.x * line_x) + (p0.y * line_y));
-                    Alloc _1354;
-                    _1354.offset = _1039.Load(32);
-                    param_23.offset = _1354.offset;
+                    Alloc _1353;
+                    _1353.offset = _1038.Load(32);
+                    param_23.offset = _1353.offset;
                     AnnotatedRef param_24 = out_ref;
                     uint param_25 = fill_mode;
                     AnnoLinGradient param_26 = anno_lin;
@@ -691,9 +691,9 @@ void comp_main()
                     anno_img.linewidth = linewidth;
                     anno_img.index = fill_img.index;
                     anno_img.offset = fill_img.offset;
-                    Alloc _1382;
-                    _1382.offset = _1039.Load(32);
-                    param_28.offset = _1382.offset;
+                    Alloc _1381;
+                    _1381.offset = _1038.Load(32);
+                    param_28.offset = _1381.offset;
                     AnnotatedRef param_29 = out_ref;
                     uint param_30 = fill_mode;
                     AnnoImage param_31 = anno_img;
@@ -711,7 +711,7 @@ void comp_main()
                 anno_begin_clip.bbox = begin_clip.bbox;
                 anno_begin_clip.linewidth = 0.0f;
                 Alloc _1410;
-                _1410.offset = _1039.Load(32);
+                _1410.offset = _1038.Load(32);
                 param_33.offset = _1410.offset;
                 AnnotatedRef param_34 = out_ref;
                 uint param_35 = 0u;
@@ -726,7 +726,7 @@ void comp_main()
                     Clip end_clip = Element_EndClip_read(param_37);
                     anno_end_clip.bbox = end_clip.bbox;
                     Alloc _1435;
-                    _1435.offset = _1039.Load(32);
+                    _1435.offset = _1038.Load(32);
                     param_38.offset = _1435.offset;
                     AnnotatedRef param_39 = out_ref;
                     AnnoEndClip param_40 = anno_end_clip;
@@ -738,7 +738,7 @@ void comp_main()
     }
 }
 
-[numthreads(512, 1, 1)]
+[numthreads(256, 1, 1)]
 void main(SPIRV_Cross_Input stage_input)
 {
     gl_WorkGroupID = stage_input.gl_WorkGroupID;
diff --git a/piet-gpu/shader/gen/draw_leaf.msl b/piet-gpu/shader/gen/draw_leaf.msl
index e20fcb2..06a4e23 100644
--- a/piet-gpu/shader/gen/draw_leaf.msl
+++ b/piet-gpu/shader/gen/draw_leaf.msl
@@ -230,7 +230,7 @@ struct ConfigBuf
     Config conf;
 };
 
-constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(512u, 1u, 1u);
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(256u, 1u, 1u);
 
 static inline __attribute__((always_inline))
 ElementTag Element_tag(thread const ElementRef& ref, const device SceneBuf& v_225)
@@ -606,9 +606,9 @@ void Annotated_EndClip_write(thread const Alloc& a, thread const AnnotatedRef& r
     AnnoEndClip_write(param_3, param_4, param_5, v_201);
 }
 
-kernel void main0(device Memory& v_201 [[buffer(0)]], const device ConfigBuf& _1039 [[buffer(1)]], const device SceneBuf& v_225 [[buffer(2)]], const device ParentBuf& _1005 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
+kernel void main0(device Memory& v_201 [[buffer(0)]], const device ConfigBuf& _1038 [[buffer(1)]], const device SceneBuf& v_225 [[buffer(2)]], const device ParentBuf& _1004 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
 {
-    threadgroup DrawMonoid sh_scratch[512];
+    threadgroup DrawMonoid sh_scratch[256];
     uint ix = gl_GlobalInvocationID.x * 8u;
     ElementRef ref = ElementRef{ ix * 36u };
     ElementRef param = ref;
@@ -630,7 +630,7 @@ kernel void main0(device Memory& v_201 [[buffer(0)]], const device ConfigBuf& _1
         local[i] = agg;
     }
     sh_scratch[gl_LocalInvocationID.x] = agg;
-    for (uint i_1 = 0u; i_1 < 9u; i_1++)
+    for (uint i_1 = 0u; i_1 < 8u; i_1++)
     {
         threadgroup_barrier(mem_flags::mem_threadgroup);
         if (gl_LocalInvocationID.x >= (1u << i_1))
@@ -647,9 +647,9 @@ kernel void main0(device Memory& v_201 [[buffer(0)]], const device ConfigBuf& _1
     DrawMonoid row = tag_monoid_identity();
     if (gl_WorkGroupID.x > 0u)
     {
-        uint _1008 = gl_WorkGroupID.x - 1u;
-        row.path_ix = _1005.parent[_1008].path_ix;
-        row.clip_ix = _1005.parent[_1008].clip_ix;
+        uint _1007 = gl_WorkGroupID.x - 1u;
+        row.path_ix = _1004.parent[_1007].path_ix;
+        row.clip_ix = _1004.parent[_1007].clip_ix;
     }
     if (gl_LocalInvocationID.x > 0u)
     {
@@ -658,8 +658,8 @@ kernel void main0(device Memory& v_201 [[buffer(0)]], const device ConfigBuf& _1
         row = combine_tag_monoid(param_10, param_11);
     }
     uint out_ix = gl_GlobalInvocationID.x * 8u;
-    uint out_base = (_1039.conf.drawmonoid_alloc.offset >> uint(2)) + (out_ix * 2u);
-    AnnotatedRef out_ref = AnnotatedRef{ _1039.conf.anno_alloc.offset + (out_ix * 40u) };
+    uint out_base = (_1038.conf.drawmonoid_alloc.offset >> uint(2)) + (out_ix * 2u);
+    AnnotatedRef out_ref = AnnotatedRef{ _1038.conf.anno_alloc.offset + (out_ix * 40u) };
     float4 mat;
     float2 translate;
     AnnoColor anno_fill;
@@ -686,7 +686,7 @@ kernel void main0(device Memory& v_201 [[buffer(0)]], const device ConfigBuf& _1
         tag_word = Element_tag(param_16, v_225).tag;
         if (((tag_word == 4u) || (tag_word == 5u)) || (tag_word == 6u))
         {
-            uint bbox_offset = (_1039.conf.bbox_alloc.offset >> uint(2)) + (6u * (m.path_ix - 1u));
+            uint bbox_offset = (_1038.conf.bbox_alloc.offset >> uint(2)) + (6u * (m.path_ix - 1u));
             float bbox_l = float(v_201.memory[bbox_offset]) - 32768.0;
             float bbox_t = float(v_201.memory[bbox_offset + 1u]) - 32768.0;
             float bbox_r = float(v_201.memory[bbox_offset + 2u]) - 32768.0;
@@ -697,7 +697,7 @@ kernel void main0(device Memory& v_201 [[buffer(0)]], const device ConfigBuf& _1
             if ((linewidth >= 0.0) || (tag_word == 5u))
             {
                 uint trans_ix = v_201.memory[bbox_offset + 5u];
-                uint t = (_1039.conf.trans_alloc.offset >> uint(2)) + (6u * trans_ix);
+                uint t = (_1038.conf.trans_alloc.offset >> uint(2)) + (6u * trans_ix);
                 mat = as_type<float4>(uint4(v_201.memory[t], v_201.memory[t + 1u], v_201.memory[t + 2u], v_201.memory[t + 3u]));
                 if (tag_word == 5u)
                 {
@@ -718,7 +718,7 @@ kernel void main0(device Memory& v_201 [[buffer(0)]], const device ConfigBuf& _1
                     anno_fill.bbox = bbox;
                     anno_fill.linewidth = linewidth;
                     anno_fill.rgba_color = fill.rgba_color;
-                    param_18.offset = _1039.conf.anno_alloc.offset;
+                    param_18.offset = _1038.conf.anno_alloc.offset;
                     AnnotatedRef param_19 = out_ref;
                     uint param_20 = fill_mode;
                     AnnoColor param_21 = anno_fill;
@@ -741,7 +741,7 @@ kernel void main0(device Memory& v_201 [[buffer(0)]], const device ConfigBuf& _1
                     anno_lin.line_x = line_x;
                     anno_lin.line_y = line_y;
                     anno_lin.line_c = -((p0.x * line_x) + (p0.y * line_y));
-                    param_23.offset = _1039.conf.anno_alloc.offset;
+                    param_23.offset = _1038.conf.anno_alloc.offset;
                     AnnotatedRef param_24 = out_ref;
                     uint param_25 = fill_mode;
                     AnnoLinGradient param_26 = anno_lin;
@@ -756,7 +756,7 @@ kernel void main0(device Memory& v_201 [[buffer(0)]], const device ConfigBuf& _1
                     anno_img.linewidth = linewidth;
                     anno_img.index = fill_img.index;
                     anno_img.offset = fill_img.offset;
-                    param_28.offset = _1039.conf.anno_alloc.offset;
+                    param_28.offset = _1038.conf.anno_alloc.offset;
                     AnnotatedRef param_29 = out_ref;
                     uint param_30 = fill_mode;
                     AnnoImage param_31 = anno_img;
@@ -773,7 +773,7 @@ kernel void main0(device Memory& v_201 [[buffer(0)]], const device ConfigBuf& _1
                 Clip begin_clip = Element_BeginClip_read(param_32, v_225);
                 anno_begin_clip.bbox = begin_clip.bbox;
                 anno_begin_clip.linewidth = 0.0;
-                param_33.offset = _1039.conf.anno_alloc.offset;
+                param_33.offset = _1038.conf.anno_alloc.offset;
                 AnnotatedRef param_34 = out_ref;
                 uint param_35 = 0u;
                 AnnoBeginClip param_36 = anno_begin_clip;
@@ -786,7 +786,7 @@ kernel void main0(device Memory& v_201 [[buffer(0)]], const device ConfigBuf& _1
                     ElementRef param_37 = this_ref;
                     Clip end_clip = Element_EndClip_read(param_37, v_225);
                     anno_end_clip.bbox = end_clip.bbox;
-                    param_38.offset = _1039.conf.anno_alloc.offset;
+                    param_38.offset = _1038.conf.anno_alloc.offset;
                     AnnotatedRef param_39 = out_ref;
                     AnnoEndClip param_40 = anno_end_clip;
                     Annotated_EndClip_write(param_38, param_39, param_40, v_201);
diff --git a/piet-gpu/shader/gen/draw_leaf.spv b/piet-gpu/shader/gen/draw_leaf.spv
index 77ed9cf929bbf122674f3cbc92726f53b8616491..7e92cbb34da93d54a2f20df565bbe42b9642d93d 100644
GIT binary patch
delta 10930
zcmZ{q3w&N>eaBCl_bt8AQUWa{P191!{niRdEo})DONBy;KufvAG)>x+T-qe8MIEO_
zV6r)6dljc!H<*Hg;ACaGjo~&05fl(nCl{euxu}5XxU^iv@AtgtpPppw{q)K6`~Cj^
z=l_3h=Xp-v^oH9<+;IDd>n_~;7qv;US28jwB}p=olGNPM^Z5Rq4<h#ttZnZ@DP#Nz
ze0^(QYj@87X9%xNw!`~+H{v@oG31)AmgT*Dmn`n<y>$JOMaC#6FQ6~#Yu&iCx2Lyr
zb%G~a<&OHgR7(w!+qyc}w{&jGynl_>+}hXPGjQ^y9oY0%^FvS2jar+$m}~a(U2w6N
zHH@mr%kbVy2TYOW7+vMlT2{68w;R|XV-TCVx_aB-CLL30omzM3SkzN`dpbJTsE$nJ
zetKHkyV|?yO@FRewoaYgpmhY+jh(9p*0yx^v<!51wPzlCBx5U&Os%gOTX|{f<SF`<
z@mqyU>tIfzR~j0sQTD*PzI9-2fBTwY9?Kf4?MLG=(AV12pWCzERkk!tCdcO+8nd2{
z!-r1jm5s@~duE!k?YaGWron!_)6mGzhH2ut)9TZ4_*O^0dm+^Pv@s*rF2wj;a}Fng
zo#?*yj?C5w@67!U2FEr!6kJ@O!-w#4vY?by9-THZT{uK9B~8Qlq9MGToDE;Ks&`XM
zZ%0Re`v83?B^QR)RbA({;MqnhCF_E-=k?zZd{wd_6E@{Sb-?B1a(Gu~PkXk=+2G2_
zm1rHEU0p5Ry{p|(?)&R1H#XKS{VWzT!E2MP;O<ttRZi|ycVW6(2dXVf$>VvHYNk^1
z1bp{2m6N9`HPh=ZdTxkXN?shsUmnI^8NzFmSK;)ZK9rKzhS<u<&dP@Abt$G~{V~-(
zj~CCb;;G=(n=Z>*P6zk5wRW{<V=E;`!JU<s>=`R1?+bpJZA-~w`0lZkk~4<!Gk0_I
zmy+h)0=%4@1J9<ovzuqJ8uZ*@J}cm@Jw3fG1)sJdKBZ*sFuraW?<#M*V#dog6Fttx
zl4pO)Y2<&QvSsG(BWorqbC)dU=dbV_2_BVsmIjZj{W8A$__B?k#cTzSH2!7nJ<1Nq
zEW7y(xOOvpGWnUA9WIrc1ARniA6Pd(jvEYl@^W@S*-*{L-vIg<Lv63nDqHdhc3n$!
zS8tCWwW*Yd;PBL@hgNwck1!^;91<2HICXj-QPAv<8e6LxwK8K?brj=s%L!p2!h~Eq
zv7p&|Bb*dkWy7qxrZ|>j1i|(6<Akf#X1J6cS;Yw-zwO>xrBO9Cl?RSoncj=nC~l>B
zOj!zcA9zVkgL`2)U6uyIx$<h$z^<r#I)1;Q_8w3(j@(sN)BZ}1qpd4!%xTF!bF$Tn
zBvI?||CY+LM>S^0wJT@%J*iFQm`yMDm{vO+el+z2xN8)AGTarGo3$S9dWQYOaMxLW
z5Oph@&*GeSCRbO5!3NjEy`v`6;C8S<b&=nl^O%q4;bMmv31;wICduZ7a@+R9jkoQ|
zqfe|!eNl0fI=q$CzNi$rFDga8eHizZp=kG&A-Gw5WhioA8G?T<NzTr$x)5yORTtb2
zdEFJcS6-1HUf`4QUr>2$c719jFS+p4;w2YcyO&(>``DRYa7FF~SLEx4aWA;Auc5v7
zTh7zOfL?InU_vjqBEN1J_lgU<0lnsePl0>Y1%IA+UUR`6px0cLr(C}XPYnvCY)^XQ
zMFS>3yYYhCL2ta^w)e&hK7kH+;{`XNw_fm%!FL8v*-iGCGjq^E?j;xv<X(iqow4b|
zxR+tk?qwL<hs?_`xC8Psl&5ZIN91i74))N?F!&r2dKm_HL|%l!*TB69gPYJxFt`c5
z1cRH;OAtOLd!E)HbZxul*hfnf@8L1|4&_#!&Nu1XJ(cdcjeB}THZezAsJ~m8F|V%s
zJ`rU;K0d791K&-#2QT^k;9a@?H8f9>2heP%?F|mlA5i--kGAhq)OK<-h$G`eW=tPO
zd597je+VvSd<1Sr^(*pT{0MA2ZAQC``ccX^DB6s59&LV{>kH_W2aP;zKLw+uyd2hM
zr;b7xe+IsmqRkUXK9>nF=wslEsXeK*>7%duE`}IGwcbug-o@8aob1Q3eUq{ysT?+M
z%yCZul_LF<U_-{(ehzkQ`k7AtWTkUn{j8@z&erC9W}XJC)v>qq)b@1cmU)fYbvQSl
zmuFB{&?P;#gXOU*JHTF*JFwdySN)fA%5C|IKMT$@Y_9n8a6hiNw&y5nuDF=kX|~B;
zf(#G-GMaiU?r*_rE_xH`y<>kzd6Duknsz(po%?%=wq^P7|2Nm{j05tKcol7XZA-PW
zBUaEc8oib`V#)>^y`DANJ&$tzT;Ld)b<Elu?&T#rQQn}$JnaIj?I5!A^cL8U^Q7%f
zikkBzj%aUZPsTeaZ&PAo{s2~UVq&KdC|g0t<y9y*%3a_QV0Vq}!oL<xU4O4id2H+w
z*o|!*Y}F@YJ4Ay6kFHurS4U#o4BwG=%I&fTn);Rbs*D1wdH9&vJ6WxXqv}EY_X7J`
z<Ldf4V?96N4l&q71RcJe2iqI0_CV%T^{Im!G!pIuR`aooCu3i@ZME46Uj)>m6Z?Uq
z6TU+|oE8#HLC_%{i2cFEgj3;azHCLp2C(h4HM1_)6K@(=-46IWwp{E>hcoCbPY1gp
zQo-vqJb-#8#g7gL<_^};p;=&qe4SbHOg;##7U#Wu2E|wg7f0Jez_yJ&au|4foQ%Hi
zD7fCf4jlph4C`&9XmBK0-I4e@B#%vX6xdB=9BumeN~InhI2!D71V0AcSYrNTJ!c~*
z7T0qQ+=klhoR9kZzy_~wpz7W@7Ow75;$}Mzd_T4O#~5}<O~2^C@!-+yzd_;sV5cZL
zbOKnd*rB;_8)}O^@By%T^(?D)avt1HdPGOU6XEL7$@$<kdT9`Qsi>okCeizoz*n&p
z+7^Q45%6TN9oE)FEq6(5cM7<k+BVwsSwO8Gu}{sO|J=BU8i4?OqeH_Ff}5z-;}Bg8
zR`+ATYv`C-bm%m&+O7EpT>@7dTbP9p4e??9Pe)MPh11P>!!zJE))v8+f{Ve=gsatZ
z;6%H#;I`G~S-gzek7u#AX37TMo_spQO=*5QodX$9aZ`ObuOsofaP4t8od;HPbH%1y
z4j;ZLSD>lK(SANy&3zJkEM1A<2sFlqx&UkvZBFyq)N0Y8i@<8JkuC<uKGGI`Enscy
z8T~;d{Rmj?6U?%FmFutAN-@!e)QE!tR)b9x9cu@x#jWk5)&Y;?+9UZIu(n9P7OWP@
zJHcyx{v&xCf?@=?1YAts1y}QNjFET4Bf0iy+XL1X$$P<Sk$gQ^tu^I`BVR|Y7y<gg
z#pL~PwKAJDk`KTmx%Oy#DOg)1{}ZrUB;N>DbKt4THvkk3;JRD}b}52?6zuE-{}|W~
z2md(O7|m>Ezmxe>urX|>uXginqSj_RZ(RnT4jJWUjLwM9Ym9OWJVvQK8eaj{_VK*a
zp8%_w=yGa#bnq&$+PCt@`_I5?SG)eV&=-62NdygUt{oHuT?to@K-Ykaf&Lt>c6IJ!
zoWD>j973Ohs};|ezl7WNQ;yby;?vZAJSa3?OHuQn5XY$h3S8V5e+^eNc+6M@9y6vr
z+I|MC&4FD<El=ak<a3bm6c3>5^Ew*-4P1L1K%WPz72l4&05@nHEPo63Tv3n5_wT^A
z)n+F?>-sB3qQ3_t4vyp>z%deS;dcXATkspfwu?mn2v#ee6JLZ!EbY~{tpArF8Y9?E
z;9{_w;cCU_{>$(Prajtz1*|RjE#P9TTj6Shk1_d^uOdV+9is7VU~N8Dw^M(O;>hi=
ziL|@>edPWLtleqdN-d8#Uk4XA<Q;G|_e|>M+XnID=F|Ak6g4-WI0F3(xHy_`z}1S6
z>o?&s812#aU%}df{~MTCX%cJp?+A*;oqQKOf@zP&{{hz4%=3IV^|vS{P<KWSW-$Mm
z>(0oXV7V{0F)-_0NjhTiG4pXVfhXX9fs4=if5X*$&YiLEQ2TMlw0)bR=8TEGmi?;d
z9<a;dc@nYjg{x;iDNf(bo!k`RbRU}fRRpp_-virFJ%Zm4Hn_Vy+C2bQ51;SnK5pLd
z`2m``u~UOT2(h7txFj9|n^fC4&PV+o2Gd8z%(p28{fD{UnEOPcN6?+c7{rgjYWlk@
z*wRGNb&2V-jsTB>4KR);zz+NvYzO?RS)0p|2g+==8;R{F=<axz>M^jNw|IG1mn;30
zAB>-`R{%Ep8MW)D!9c#nJ_%l6A<z@la$is#>CeHZQmbo!oLbGcK9Wy^H&I*)zv$c!
zev0A=tH1o|lpitLXTXj&*6tTz8^qTva(%RqCddx3NsGUT{1UDnKF@-U;Zbc;xjx<%
z&r_!_@MAp1Ad{)>%oyrlQMARY_pfsex541QL3@r8Ba-W941W=Q2|S5n;uopq+MMm*
zg0JKcdztZNh=RuT7=H)0*EW2GT0WOAceziJ--EYOo7kYX)rUP)Nzd3feI2~OLfhA<
z<t~|f<_+)?YIW_eQmdKQCEEoqF4>##os<}$Tp#`2k8go5r^JfA?S{5NEEz)P6wC1r
z+@!@N`vY7(eBK2cBbH3APjs?|-ygQZ&*Gu*pc?^Kvv9$8G>=bffu^w=#iOT$uI*ap
z!=`1h&DA6PNN_Ry9&j~_^Uq_I?H-S}QD~8;a0;c#Xbc)7(HO9S)VJix$AayEHj~?-
zao{`Yc=+rI*H7IAqiLcRot*$Kc6KlLgfe~;Dy_|R-4mhq&#5fV!`|rHqI-2<o2p0X
zeZa-g`@+>M!{=cV+N+gkn(NZZAX!A-4{S^I7}ON78MVd0_XigTUJut#-GTcz1-0;R
z0ILO`24-^7?EH2n8xeHUU;zIJq2?}g=?(y^6{lq;Ts?dad=H;lXzFpM9|X3YdQ*~C
zgC7hrcwtW@hrrb%_@VFNa~PU>1V0>XJ9THzXYL5FNws}4*Nz1H7r-OPpN)1pwPGYY
z8a#Zo$DpY@g$AAtwp}Ej1J*~~8T8LIY7z5T@bE#V$02HnpvQxYLEjJ8N8O<Qtw$~V
zPXL=JlFS8r7?jy^`pfkR{||tTk^85~Jctc$$Q#IYihw7A6H0O2=fl+_=mM~vh?!WZ
zjiMfLP68XJxWr9-0iS4yKqrIkT?9G>+(c2|LTjge5!eLU{EM2}sbKX;_Cc_Nx-B1+
zT%WkbP6KC`82N|XVoNY+j9ctOU^PDm_ZmGNzL*k+yj(vA;Vrfl9JiQ%dXsA#iR~;f
zw^+6`XM#^v(D+e|%fR;9b!!I8;}$y`{3U!$Y)~8P<D1Z9m8LW6((_<(kDUvadynzI
zYm(*QrPS(1KZjb)%-&<?gEx`cA^PgE625{GLzL^Izw><oxOk6U=t*LOxW^DOr<k*g
z;3h5JV;95K!>0vojJU_-`e<KBr#=EUhB@qr|IecpuAZGFI<0~jB=)~tr&#&b;P-lc
zX-C%<3*P}&^AjCh18<|mtjP7VGuOM5Iw7(0Yr%4juKXq7Sow8exwZ~$-C#TI(sWVF
zW8r(iU&Y77#<Z<I*}}7AzJdCRI!9qyE^NRd68C|Nzs(K6)%weozU6aj7F2Fp{^=RM
z7{?TC!T5issB-R#(`zPGZd`E}EiO8L;}l;8zwn-(FP*RF+?AJ1iSAg(Bz}=LlPY(w
zJUFc{s5h&^V`}9$Q5@7vbVq3&gZeVIM)+p-hO_$>xLWK#ZECik4z~RP)bH@?zzLM&
Mlk&E<3+B)JBTmls{r~^~

delta 10916
zcmZ{q3!GPFna9sC|34z2fTNCFh8aY>;ssNakjtQkWGJY}MR6Kn1}5f$!?;Mbd%$AN
z-PLmct#0k=>Xw$KZZ2!>uB*FSnrRuC*~LprR7!<rtrs+Mzu)se&v0heJs)4*-|zQ4
z@AJI3^S)<hyyfm8x7<DC)0ZB!vo=XaCqt7`k|bj&$&h=ypB~roUF4p=wQap9WsE<B
zuW#*b?aKL&2k_ct2fVjuBfdiuL$2v;S<=(HZeDNC73=4pWsGvN3;nF#){P5#x_dfS
zCwQV&?j2v3YN;Xes?Lt}Ege^8-oHj$*xK9H-8c7&c5Hg9`Op(|qt+(7bIo4*U$K`%
z2JxZro-6uHk>wa(<r7<0wr*%Mut7#3Hg|URtb&_#M5T2?-7zCk&+O@L?^vTcG?Dvs
zx3qP(b<vv*xn8+xLVbhQAy_wdtnORe($U@0*U{OQdF+>rtUNJce9g$p?g{l{^(*7I
z5|`HgoJ6lQG*qMPhjo2x-`WjrYX*5NYN)o~ACJD?*6t0tJ?mX%OG7<5KHJck^?W2g
zbV9FeOy)fx(~ND;?bnV5`?a^Bk$)Q|il<K;pB{j3b>#aNLd|CtGh*#RjLJ3VupaD0
z_qMfXwkCK-?sqgew$ZWR;(|>bz{|<(Qc`(x;+S;K0KJsV9mLNbz{|;n@Rciju59UP
zZ{N_?M;}VbiqJZ%>%2O6wvkH7y5Q`2{d<G2OlD=mrd+5FxSXW$&W`T3Y>~6Um6NS#
z?H!$+EnPjU-BIrQ>npc5)-AXJ3z^`x$<5%dR=ZVB?o)SRI$QgyElSBVd6a6VQt~W(
z-!zqz=PNZ$bt_&Rpq7%|gZL|h_^ShWZSoqN{?mt2^7;T<IoVUWys0j&VQ;QKsoLjJ
z;%Qaf0A79NrmW>8@P<{boo(6JO38_EXQd^3$4bd5!8h5ql$-<KH<nVeU=TllA5XIe
zi}nfda<UknO>svT?_xFR#e;m7!CSk#ds+%Us|Wa$l8!-q-5}mo-gecbmutp&oQ)vQ
zILbNX-&NT%dEb#W8I`$9=JD^Z@Ei&rmU))?kE_EnzVG<5jn6!`f=3$vW$iu64$mz6
z_$0V?GkY@mOwJCM%C%GK%gQ4w(~hW{8AlGryaQNhU}H6pfXo}sIBKIqt8B?5*nKV0
zWg~b5sZF3n1c#{B6k6qpJi_qYa&%aT;PmM|wV>G_HMUnZYGu-t>L?D(EvJWt2%~cC
zjDlwGjW8#)%H>n)n&W7S5d_!M$4OVK&2T9@wu+NJYTLt8O2cY~R31BidHN7up0CF-
zWlqgE;lsG8X24Uhd&7%sBHT+%A4>z_-3@9J!LG8r3BTV^dpD>VNAAk2X@51xw4IwZ
zcG(R&F*!I}(MS@t_W!q3o;sm1JHlN_!yia(D(ArTa*u8`hj}9PXt-+@{4lsH9@8)Z
z?)rxPBDm`=KazSSoM&FnmnBzMh5iON!kyQ88r%bRD7wfW%z4bmOK`D6j0A`Ct|rOm
z+1s8y@r;_(j}<qm!&y%4$4Zg=u~Ota260mt?S3`{x0il46uF-b!Ea1hC9k^>Y|vET
zcE~F)>~gQYBA;5>HEn!ak7Z85Vk9rR@YL=_7hJm+UGPWPiC%I=?j={`>jrTzx!j&6
zYY5;aS8U)V7u<wia7BLoAnr96cH?=~1s@Cdx(ogy@x1D)oJ@|;yDl6qhd*E7?m=(8
zurvADtry%5dg}$Zy|-TQQFOptFSz-<`7)le4L*vnCj_`X_7aQ+axcOn_c9Fb5_lmN
zxff!Qdm#pQKwgOQ)a~ntyb;5}5qTp9w}W1Y!A~Zk7h-Tn<YgFq4cyBxxCy-ogPYHb
z5I!z@e@y627!GUTUE8iX<;l{Rhk0$jLGe(0i@rWw>6+enfJb68x6l^qM=FzM)OCGJ
zM45??FZQG02PhBWC4UUOC)e*m^W^w8n(ef`!GZZWwU2qUeTSm<CP#rdGHy3x`aP8G
zl*sr5xR~+#a5Ji3mG|OFu<f)N?SASXP`*mhW~_^8^Hi?SqE{X?^055~jF$3&Seu<X
z3Ss;S_&SO<PayepCcvOigD<1@q|&C3zUq4zVhq*#1v>H$zLw%-KZEUSl$}ZC*cro5
zeHN$`>7N4|GRF2(uw&EDbn@pa9W%yHc^>3!ZO&)r1+ZEjdrMDkFI4WB(U_WkdOj~d
z123gZdU&A9V^w|*_G;XT-Tt`hzksXXm9O|O!D>6Px#BN@eOz&EFH+Q8aWSzgjWf&X
z%izJWUqMrkwf!wv%{6Z(yI1P(D7z`!(X`tyujk)Wv@Oa<|3A5AUmTAw#A|4dk+#dT
z0W0Vjja~;Eh#Bi|^ak8UF^zKlT;CX&bxhiuZsqxVP~N1(G`(dD%1$CXO@9FUI8EB#
zrl>hh;)u34doliqvX>GQ^A1?eiHVI;!+B>vT}Xl^a--Z2tp&SjY#06|G<E&GDdn-R
z%V77laj;ci41;TM;L%m<=<0shHp6%3opQGfLsS1mzAF2J)$Zg=z{Fn3YRw!}kK;cE
z>{pDd>(`6*On*W%*g*(7d_50V2UdG5bE^6r3^!;bJOr%fTNf|Jp>W%3vlD&?s6{8n
zf}<0DGHp)_3C1Dl5bwiya4}&6T+MH*NH`H}J8cV@(;JA_2v)ZPes{>lP1Meyvpfmx
zhDZgk({M8N5fmRCrsNLR(V-*32KfrJ<dJ+7Sj{PSx5*Eu7|Y<|XnPFUwy{T!1GDVa
z_4jK>!S(j*=XmgESZ^CegA>5&j$|seJT}#dU^kU<wCUp)m3nmGB(Td7d>Xi^#Qev4
zo{XSaT+jEwZK%!8`Kq4+Hh6UdRrkiJaCMIoH`{68N2%RE#;`+b`b7uc3m(q?>laQ3
zJ4Ml<>0q^Dhu#Obp|;oqGr;QAv#i?5GvId8qa_l~gsVp<XMxk`r9teaqK-B?l-@Uk
zx3U!4=78l9@Jz5B);5<~?vmK<Ebw@0+i25gHnn=he!uU(qQ<kS5eQI*^P>T{nOZ##
z(R0A+KDN1rj;Te5=7ZJl%s1!<;A$fZvvBSJAJ+ea2#R}fx;<~W0B&P#5&S%GG5Gm#
zwOV$6v|9+btv1i%3#fg3JG3pL^z!ZG+aYdF^V4ZDWE90sbx~eN;)~(h<8WF6R&#U3
zrd$diyeXHVsmIa29IWO(i9MEHg5U@=#)i5SY!Yow^M%xE(V@%0YO#@8z_E|Ch2Mw3
z+SW7rBT3o{R{J=!EMMjNE3TxNXazN5e}FcyiK1ieV70ilebv^$Bf0iSz80)4l6QdB
zBKe2G9lrmOd^Lh%1n2}8lXt<@d>v!t-S9}RJ=*qwwMFvvV6{koxi-qmlz$xgI%>rT
zumN05-UnAJvq>ZQ74S%|J=*>WSX(6D2v&>aSAx|Xcxv)qfT96hmydv5ir^mwJ3GNY
z2DZb&{}gPDg=}VjCG%%sW7tk#?dI7;t<84cx|?19;wZOZM9kjT80A&)7^U`T{Bf|h
zk0sSke*&y#B4<?|9sF~!+JpJ){YkLe)vo^?^u^v>gP_69wUc6?t#I`S^cUb_pudEx
zU7h<F=dZxUbK+XKTJe0j4sP3P4dy}dX=)!23XPwlsCiI`W7K~QF3wm5u4eF<vCqI`
z#<WM<zX5A=VAoU2)A%y^TgWJi2hfdq9SuJV*B%GZ=fG;kucN<%8#E4<zXyA+sK@Jj
z6WF%e?8FVOzhWf%2QXs)NNxegNVJ9Dtzd1z{}F7vNc4HITJfCt0z6`AueN3VZ-Z!z
zV7G&d!M+GrE57$%f=4jz(e@6ow%~Vyi?O~8SL=U`$)DVX5W#ea#&?6Y`C8pW{ZAA}
zZih{z-QDji_Z6^qr}buPdBnLFT-=cV3|Di{q;9@{f%v%jG;X7)x%tEq=&Ru3Xubwl
z^LmZV_pk66jP_{zZ(wb~{~b)MG>J94A3?FWlm7!A!L&!?2f*4E@;*OE{ht&Qs5>J^
zF_{0#b!X&0u-p&Z7?|~*BpuTKn)$k!z!UJl!6_vUjIV>$e9xV+hp2s=F>T+Vs5xU|
zuVsJJ^Dx-u@H~mw--N4YJ}FL*<W6piaQYUS`c?wjp+~_sRFB|~fer31k9OaNtB22b
zavwKu_&koLZtRH$|1QLa8sd`J4t}&$x#!$Norgp^l4DZ$C%|+yX{3A)9C5!7);aq8
zBv?&<S3#Sas}PfA9qoSr)@LO9-8@f$$#YTmi)4+?egx8J!=X6-5bOqbWu6B6yv?_|
zx?Jf;{KJP=@?&b+RA!%dNctRjHpQ6FQp^4PbYMRPpG~c<{TXUCbNa%)0NzA#qWvxA
z4)F67&rkj3`o!RV26k|<FgxA*Hi#eo5fYtr*q<Oj2b;9`W5zGw>f!TCurWM{O)A&N
zm+>X)U6fH2V;n|pXGT!}in7bQFFve)os0PO2mcM)i<B6VTt9>OW9G}?dWwm6Q_Hov
zBEJQ1&26tx%eAe?_B*hho^=upUj-;k=dTCc7Rm3yH&L6|pf=XWJ@a&B+TyzO4cKgo
zX<w(7yJ-Hk;Z5*-YIW_eQLCBSMSBZeT(r00dnhqPxjy>41OEU{DY0mKT{IiSq9J5X
zu^@kho3yxS@4(f=hrVVpV$tOKL??&9Z-$@GYv4gv3s<wSaoZOjn3jO1u^YvsrHroa
zI%3*%DA?xe5q>{#G5j#Nn#C31J<9fv$JhR7k*9FVq{(m$8Y9sNuz}RK<jF^Z?SM9u
z+o1!%_tEk2IS{U&x(nuS($%7~qrt__j)9LZ)5X!1)`fLl2SM$hQ(2scI&^K(y@SCv
zRgchzfQz9Ig{xTx&qF=hYn2@f>(awOvWPquY)kbR)HtviwZ*{4gNp;70M}35f%_)|
zweX(^Rtw$;W^&T(e0C<A5OmUD0RP6I*39g?bd$kq#c4SLt{y&9-o@ugH1#;dj{@6H
zy*Wv%!H<R*ys#&dW8mr${MdK#ISx%df=>n8PJIg<_nkW)Y*KBX%(WB1hf|Ixe>PhG
zf}$A7P67`e?KCuXr_jJBgKZbd-vic1-5K<+FlrI=RPf+Irl%olh@kHU7lWP-*GJu;
z{wYT-{HKFW6iMC(_CP7K<@A^96aF*6#>oBC<P3-nZpj<Sb&7y9!EBTP>plyv9zkb=
z?L^GP9I!s>5vLh!oZ=GCEpugPh(Kq8?Og;q3p|&iz9nC~v%x0N<{!{piuZ%nBiTH#
zgSsmplw6;<$L52xdyM=8F0v0`&=~jFxnMOPgL{#F5Pl9N4tcqL4#IovJaF7&{$Wk7
zZ78;d;JC-m2d4@eKZ5ZBu)TKO7J=n)k6j4<0zM`-sEzf>?lIc<q4*POWBKc{CD`It
zyBI9@R^xxCNtS{aP@`9WUAh=S(Z*hD%fXun=7{|kxdgt95+j!DqrcOCDY$s8t$^zv
zx0+m^n8C}y=G;#HEQnLnf}kOsJ_I&MTx)WDw9lqftzd(g#E$rXb*zM|htDdoG2(5I
z>l6E+Erq<>2TePA8zuI^8nBvAbZ{+vH6><6uAiN8AAA@b`=A3X*XCe5!R~|X*Cp#9
z3L4unc7yG;`=E<j9{ZpN{AGMhY*5?kQ``rBqRKCzc3<>jkHj0mg`eF<Bz*{q57kzB
zmz-QPyK>u-PfhZ}IOb>z#{V-#mBmZXt*Nivy7YcptXQ^jtY3$pdsolh%k-SSeBIdS
zk9ADq=V?=4d0_cb>4buMyDB`UR(>1BQ5}KqsH|gDU&Ph~-^|`{dcOo$iw&qv&GwVP
Xwx3MBm!Ahtr<|6Qx2?Kl=8XRZo*woC

diff --git a/piet-gpu/shader/gen/draw_reduce.dxil b/piet-gpu/shader/gen/draw_reduce.dxil
index c001e896f2b6b3aad03bd517ce94c82d200bb760..9b1b0fd6893fc7bdecd4093fd10ad0e5e93d2257 100644
GIT binary patch
delta 378
zcmaDN_e4&_CBn&hCGYK`+5hwdqBdIJZU`z;oG8-C$hdKyD<dP*<Ov+jo70(88Cg66
z8Wbi^VD4p3Dk|Enz>>i%ptR_?wL{5=&u32gn5nRQ>6<)*{W~Mm<_R1d7#Te_J97Cm
z3odeC_A*>h!F<YmK?yV0j0F=WFXYvba{SaQA;G}W)L6*IDbWLCdP;;$nEaSGl+kgr
zC0{^&5(C?29uC1Z1}#mp5!yTp3)ru3U{sV~Hpp>cIm&QGNQGIHp*iKW1M2~Wn??bQ
z7nmpMDzmXMG$k}N9FpAO_|dUIw7}qwz#T@(Ge#QT8xBq2s#92Fu;h?Dk4l3hlZ2{9
zKoiqu7NC9yjR>HA0j_BfmoRSPDw@E4IEIZyh3CKq#*Zw_31>79B<xUJz-%xhHDOZ1
z4o2PpMxHcDDRm7~jaF&H3<d!f<`tr^7Vt9J8~i!I|K|a(mPV^AMEr#4n*+Q}_6GkR
P@c(1rOXHMbVE_UEkx_b|

delta 378
zcmaDN_e4&_CBn&hO5K@?n1UK_nZjvBcS}XzO%&;5WZF2-m64HU@&u0N&FM_4j4T}i
z4GxnhF!wU2BqVHBV98(>aCmk4eA`FCjvK4|d<8a64VXNG{W~Me<_R1d7#TY@J97Cm
z3z{qtVeT?tP{ZtHzMz4VYo)^Eg}fS4jh}iYBp5iF8VlLDBzj;>FNu%^lOOYjGB!@O
z<O`^mmN;Xi;IiS609RSUT7#BDvN4C*S`v5;Y+w`=VK(rdz!Sp2Gf9_Ol%YB0v;(UH
zzsEL+NsLTf84uWwvL6gRaInC+fIY?F^a7oOT+UiQ8aEu8z*TQBfq{)PVafsrPOdXT
zIzask6HWs4YeX0&B`_#(u$|C4<0!`?9V0NofZ;IP5n-MN!{rRk58E904<t-$^lW;_
zX5zqVFe5c%&VnloMDmU-U~~}RSrE3N*^J}Qfe*~~ADYcqED$M!h%1EeU^eIYci=~}
P{SRh~8L9;W3?KjiTV{gD

diff --git a/piet-gpu/shader/gen/draw_reduce.hlsl b/piet-gpu/shader/gen/draw_reduce.hlsl
index e56ec3d..a6ccde9 100644
--- a/piet-gpu/shader/gen/draw_reduce.hlsl
+++ b/piet-gpu/shader/gen/draw_reduce.hlsl
@@ -15,7 +15,7 @@ struct DrawMonoid
     uint clip_ix;
 };
 
-static const uint3 gl_WorkGroupSize = uint3(512u, 1u, 1u);
+static const uint3 gl_WorkGroupSize = uint3(256u, 1u, 1u);
 
 struct Alloc
 {
@@ -49,9 +49,9 @@ static const DrawMonoid _89 = { 1u, 1u };
 static const DrawMonoid _91 = { 0u, 0u };
 
 ByteAddressBuffer _46 : register(t2, space0);
-RWByteAddressBuffer _200 : register(u3, space0);
-RWByteAddressBuffer _214 : register(u0, space0);
-ByteAddressBuffer _220 : register(t1, space0);
+RWByteAddressBuffer _199 : register(u3, space0);
+RWByteAddressBuffer _213 : register(u0, space0);
+ByteAddressBuffer _219 : register(t1, space0);
 
 static uint3 gl_WorkGroupID;
 static uint3 gl_LocalInvocationID;
@@ -63,7 +63,7 @@ struct SPIRV_Cross_Input
     uint3 gl_GlobalInvocationID : SV_DispatchThreadID;
 };
 
-groupshared DrawMonoid sh_scratch[512];
+groupshared DrawMonoid sh_scratch[256];
 
 ElementTag Element_tag(ElementRef ref)
 {
@@ -129,10 +129,10 @@ void comp_main()
         agg = combine_tag_monoid(param_6, param_7);
     }
     sh_scratch[gl_LocalInvocationID.x] = agg;
-    for (uint i_1 = 0u; i_1 < 9u; i_1++)
+    for (uint i_1 = 0u; i_1 < 8u; i_1++)
     {
         GroupMemoryBarrierWithGroupSync();
-        if ((gl_LocalInvocationID.x + (1u << i_1)) < 512u)
+        if ((gl_LocalInvocationID.x + (1u << i_1)) < 256u)
         {
             DrawMonoid other = sh_scratch[gl_LocalInvocationID.x + (1u << i_1)];
             DrawMonoid param_8 = agg;
@@ -144,12 +144,12 @@ void comp_main()
     }
     if (gl_LocalInvocationID.x == 0u)
     {
-        _200.Store(gl_WorkGroupID.x * 8 + 0, agg.path_ix);
-        _200.Store(gl_WorkGroupID.x * 8 + 4, agg.clip_ix);
+        _199.Store(gl_WorkGroupID.x * 8 + 0, agg.path_ix);
+        _199.Store(gl_WorkGroupID.x * 8 + 4, agg.clip_ix);
     }
 }
 
-[numthreads(512, 1, 1)]
+[numthreads(256, 1, 1)]
 void main(SPIRV_Cross_Input stage_input)
 {
     gl_WorkGroupID = stage_input.gl_WorkGroupID;
diff --git a/piet-gpu/shader/gen/draw_reduce.msl b/piet-gpu/shader/gen/draw_reduce.msl
index 26a8793..8a87f15 100644
--- a/piet-gpu/shader/gen/draw_reduce.msl
+++ b/piet-gpu/shader/gen/draw_reduce.msl
@@ -45,7 +45,7 @@ struct Memory
     uint memory[1];
 };
 
-constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(512u, 1u, 1u);
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(256u, 1u, 1u);
 
 struct Alloc
 {
@@ -124,9 +124,9 @@ DrawMonoid combine_tag_monoid(thread const DrawMonoid& a, thread const DrawMonoi
     return c;
 }
 
-kernel void main0(const device SceneBuf& v_46 [[buffer(2)]], device OutBuf& _200 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
+kernel void main0(const device SceneBuf& v_46 [[buffer(2)]], device OutBuf& _199 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
 {
-    threadgroup DrawMonoid sh_scratch[512];
+    threadgroup DrawMonoid sh_scratch[256];
     uint ix = gl_GlobalInvocationID.x * 8u;
     ElementRef ref = ElementRef{ ix * 36u };
     ElementRef param = ref;
@@ -145,10 +145,10 @@ kernel void main0(const device SceneBuf& v_46 [[buffer(2)]], device OutBuf& _200
         agg = combine_tag_monoid(param_6, param_7);
     }
     sh_scratch[gl_LocalInvocationID.x] = agg;
-    for (uint i_1 = 0u; i_1 < 9u; i_1++)
+    for (uint i_1 = 0u; i_1 < 8u; i_1++)
     {
         threadgroup_barrier(mem_flags::mem_threadgroup);
-        if ((gl_LocalInvocationID.x + (1u << i_1)) < 512u)
+        if ((gl_LocalInvocationID.x + (1u << i_1)) < 256u)
         {
             DrawMonoid other = sh_scratch[gl_LocalInvocationID.x + (1u << i_1)];
             DrawMonoid param_8 = agg;
@@ -160,8 +160,8 @@ kernel void main0(const device SceneBuf& v_46 [[buffer(2)]], device OutBuf& _200
     }
     if (gl_LocalInvocationID.x == 0u)
     {
-        _200.outbuf[gl_WorkGroupID.x].path_ix = agg.path_ix;
-        _200.outbuf[gl_WorkGroupID.x].clip_ix = agg.clip_ix;
+        _199.outbuf[gl_WorkGroupID.x].path_ix = agg.path_ix;
+        _199.outbuf[gl_WorkGroupID.x].clip_ix = agg.clip_ix;
     }
 }
 
diff --git a/piet-gpu/shader/gen/draw_reduce.spv b/piet-gpu/shader/gen/draw_reduce.spv
index 73430c6a938744794d28695d8f82174c31aca249..aabdf6b29f9fa76cb0dff6601013851140091547 100644
GIT binary patch
delta 2394
zcmZ{kOKe+J5Qfi<b8j5CNt-&zgpd~%5*1y5P$duqQ78l?qRI=3#KyR(lUVqXI*tkw
zQrNRVa2AL+DU?SkP-DWQ&{9ZC%cDRs5SOrk9h+|0@L2JG*Y_Y(#FhTc%s+Ex=FA+Y
zx$X3}o><4`E@M2CFmdzKqdhw}oi@g7G=b@q>=TSvD?J;#^)+eb>TIqe>C0J{oXb|S
z#mHX%m+hMs**hxPeUs%<Ilsp!p*WlFOcwEY!c-wYm&xBBW;-c+=R)=7h3Qr#HeY$6
zx*N_nF>|U#Ne>K7&J;3tmn-*8RLTo;ckF0&)et|KE0!w{kS8z=d$uz*087PUu9(SH
zDrN3KKMm{o$)WFM@o}qF$?b(gc}kLqReQCc+|rbFOSv?ipJ`>q<|}1#g<LULs?JBe
zE&8b;ydA>6{GRG;CSS@_^Mza((Pf%;$shEZ_Nbp4AkG(dD<u0*%|J2dqbypCW?Y?}
z&&^yAF&4DWcPpY=$(H6LKX-20bdXYwd@vGrzFskEf~fFL)a?m5a??3|)dKoD6O8EZ
z`B3lf8r*VDajlz&s|nA;@l9zWNOrSkC%Yc@R_xw*@^ZJ|&GPYxCiM4&WFW+a+jZik
z?TaTT*>X%Eg6&==@Hu--y7RN~b;2c)|655iQh|&OYI_aa*PTqWZ#F;&RGd&2IG2l!
zx7Wg^79X{LkfhdtK<-s?+h61TnvL?0yA8Yi#K4gsj)^;t{@^w`&Ss+wbdTp^=aFRN
zaF@1m+yHLWakds^NJO41oJW$riQd}A(GT2Jj<a5(b&KA}bFK48Ziu)W+Bmu=!%dE}
z=|r;6y##lIb8K}8HYPC1HOcl#uhkDE?(Fb9``h|^7MJ8?O;;+AVeA*;o-xC+SER#%
zA*JI1!3|62*<g<<?o&x}z+E5NJS8~ntRzoIX~Xz0CGjcOVwl*yqha2V9Ihxl1-^?8
z;3FY6?}a2=6@9NH@8uEUt_U>Ae-d&HoBr@^xV%oTnn4o=5%?e2h2D`d2FbIcjUOc`
z2}c{)YxU*cb>0~NnRz$R7fIGHu$Sw<Y_PHM=S4p!ObdT06#FURY2gp$M*ms-k|fwK
zY}ZhxBww(#WODIEG+|s(T;qW7gm6Mp*e?l;0REwfds#Z?@EsQSpyVq8JR|#EqdF+b
zM_2e>72xO!J?f&v917cbUDP4Lweg1Z_BP&>O&jp8jl<G8hmUlROV)(P1bpP`mwZcr
zzfo;5S8V5dTl!v}4Nis(;hgVC=K#KCN$l%16_UOy3<#XTM~o7M{GYm`yDRelHt?Q&
zR@i0ikz``=eW_5c{1NG-ba(qeI(H+Ea@fRB?ioo&i{w13`%qwpZeB+t?lwuz<AdY;
zlg4Zc6Yz4}3pp09<|8r3goDwhk4uN!C9vsFB>C9%amkMbINCvX**=rbua<MJ#m}X~
zyDNVoy>;cl@T)$<YEDTR6$nm8Pw*@01oKhg8p*E(9SCE&A2xCB%s0|G;|cD}S?Qc{
SGr=Z~nY1nulT!Cnjo9BCS8iYc

delta 2414
zcmZ{k+iz4=6vp?N%jtBWrRAWtVs8=?i9BcwiAHNg(Q0Bt=0Qw+8HX~o6X#N<)1Zk-
zdH2ESelXrCikAvXDON!dFSX(YwK8gpHSr(d#l(1d^81}Rn~Wr!<m+1NTWhbi_F8*-
zX6y3Sws>q)i!q)_o0R$O>9!r~zcR*j7~do$I|UP)Ex#PU=diR&Z7g4vOv-s&ay(bf
zl_Gooe|FL=$lhMf?HjI?D}~)g3B}n=A}HbUgpp!lJX?4y%(f_d$7JoU$vuroY`)55
zZ5Le9#LbxoCEYhPIa<s<RH;5XRIN;o-@m=lRbBjWzEr6`P9ERX?NXwv8<vWtd?}l+
zRx8|re(Kgs20cH^;^S7!lG}>K%7`Qp%l4*Zuz5w+-Iem5!e}EaHeWfLFXl`6a&02w
zZ7JE+gEuDZE9|a~WeeqOtx(K|5iMrLPA4~dEA~{ftDCr_uuCDii#2`4T!^x0F`990
zY$88;NyLEPINz#>S~XXmi2U4n#b*32)hPIVVdrZVqb7(7??m0UkRvyp)7Pv|UuXS3
z{k;(Ctt}fjpI2Pt=HY6>^Kg74nh28JtlHt0C%px`Hx*px_Pa^GgfOJPUr72wO1Mua
z7HwxL7-q|Hfe5yHt-$B(0qM@q#@7f}NB(n?WTXNa>(ut@wzD<Jux~a%2UHwV7C4uS
zjW^fArWPNyf03kCw?OV?a@$(E!|qQdG#}L%bX#^ch=ZdB98-53J;HT3&gP;F+yKv2
z&Lhdz;jU@oxCLCl<LqpdAsu<HcOFT4Cwfa0M?Y{kI?j6O#x2?-&n?a)xh~>vZ{p~l
z40k%tX3{~YdkyXe=h<uzY+PWPtCG#rUY$FXejw(>>>q0%o;n~WkN5@!GK_;ltjT^=
zdP=|DUio=g#DE^9=TX7+N?(-3?pNHGlH`E9EwXuDaM(FXo|Do}g$K??PLh$>y{%r}
zl{i-vo&rBb@%TuH%^RY?7DYcI$s2k~crXG@^52DAy=FuBR(xIiVsab=ToyLp-Z#Kt
zd1kcnt0X1iXaifrZC~RJ9G1#Eg1$oX_$BuI+|oK5ADkBbif};qSE1N12rmkMD>wS@
z;txs!wy{k^8IgR=P6xr%Av9r7QQY8l;W=SJQP^(?i~xRr#Jwq<bNG&mdqVOp0iKcl
zrBS^t$wycCW&}98LXWzLnImBv?}$1gxHjIE-rUAf*|Y)g+L)EjIeeshR`QtejDU|^
z8zkoh_&d}VbH#SP_oVOT8R2Be5YG9&bPnK~m&CqRQz7XGLbt#fe8ebG$p51|y1ODT
zw}BJ#Sz)KKGeLUlq*N$Z{zK`cba(qmI(H+Ea@fRB?j=b^i{w13J0&nfH?NN)?p{gG
z<AdY;^Tw<X6Yz4}3ppLG<`Xfeg_&s6pGt?@DX{6!CHdGizRv_W+Cg{OzLw6fmUFJf
zZ=}P!D}O7!ab@4|qyCQ7@YC-X2+l-L@O$Y5^HJbt$sYtA2xGY)HgWFEl61~^f;;o0
Xbk4Y$U=zno8kdMmnfvp~bMgNGU_)_?

diff --git a/piet-gpu/shader/gen/draw_root.dxil b/piet-gpu/shader/gen/draw_root.dxil
index da5cfe26ab31f97bb4faefb8ee3c103d3d6faa70..a84fd4aa668225d7ed426d81a25bc91b30599fa8 100644
GIT binary patch
delta 185
zcmV;q07n1l9_StvL|8&YS{$*JW+4ZU`jhD{;=Jj^krasm0kNJf0RaND%mIY~2uM(X
z0R>TjAhTHm^8*m0ygUb`#0y!|rNiJn>a{eJTnG070+Y-EZL`4$AOQhLv*HOz0|pEV
zqCf*mlQ0c30nd|T4O9WRlei5<0cMl_4Ojy329sD0JtlfEY(PN401O)tk1zn{r~(Xz
n8-hLngINY+2u2JH00hARfrUQcle-R50q~O!4>bW|lSB_B8pl9e

delta 185
zcmV;q07n1l9_StvL|8&Y+9QzZvln3>ZTQ_AhWvThkrasm0<oSg0RaTF%mIY~2#8RC
z0R>TjK(ko`^8*mE2Q?7I>IEX9$8T<1ld#y6TnG071e44GZL`4$AOQi0v*HOz0|q3@
zpg;o(lQ0c30nU?S4O9WPlei5<0c4Z@4Ojvo2a{M1JtloHY(PN401O)tk1zn{r~(Xz
n8v;K7gINY+2u2JH00hARfrUQ!le-R50pODk4>bW^lSB_B{C_`Z

diff --git a/piet-gpu/shader/gen/draw_root.hlsl b/piet-gpu/shader/gen/draw_root.hlsl
index ec75d5c..56b513f 100644
--- a/piet-gpu/shader/gen/draw_root.hlsl
+++ b/piet-gpu/shader/gen/draw_root.hlsl
@@ -4,7 +4,7 @@ struct DrawMonoid
     uint clip_ix;
 };
 
-static const uint3 gl_WorkGroupSize = uint3(512u, 1u, 1u);
+static const uint3 gl_WorkGroupSize = uint3(256u, 1u, 1u);
 
 static const DrawMonoid _18 = { 0u, 0u };
 
@@ -18,7 +18,7 @@ struct SPIRV_Cross_Input
     uint3 gl_GlobalInvocationID : SV_DispatchThreadID;
 };
 
-groupshared DrawMonoid sh_scratch[512];
+groupshared DrawMonoid sh_scratch[256];
 
 DrawMonoid combine_tag_monoid(DrawMonoid a, DrawMonoid b)
 {
@@ -55,7 +55,7 @@ void comp_main()
     }
     DrawMonoid agg = local[7];
     sh_scratch[gl_LocalInvocationID.x] = agg;
-    for (uint i_1 = 0u; i_1 < 9u; i_1++)
+    for (uint i_1 = 0u; i_1 < 8u; i_1++)
     {
         GroupMemoryBarrierWithGroupSync();
         if (gl_LocalInvocationID.x >= (1u << i_1))
@@ -79,13 +79,13 @@ void comp_main()
         DrawMonoid param_4 = row;
         DrawMonoid param_5 = local[i_2];
         DrawMonoid m = combine_tag_monoid(param_4, param_5);
-        uint _178 = ix + i_2;
-        _57.Store(_178 * 8 + 0, m.path_ix);
-        _57.Store(_178 * 8 + 4, m.clip_ix);
+        uint _177 = ix + i_2;
+        _57.Store(_177 * 8 + 0, m.path_ix);
+        _57.Store(_177 * 8 + 4, m.clip_ix);
     }
 }
 
-[numthreads(512, 1, 1)]
+[numthreads(256, 1, 1)]
 void main(SPIRV_Cross_Input stage_input)
 {
     gl_LocalInvocationID = stage_input.gl_LocalInvocationID;
diff --git a/piet-gpu/shader/gen/draw_root.msl b/piet-gpu/shader/gen/draw_root.msl
index 2ed7ba2..0d22e4b 100644
--- a/piet-gpu/shader/gen/draw_root.msl
+++ b/piet-gpu/shader/gen/draw_root.msl
@@ -61,7 +61,7 @@ struct DataBuf
     DrawMonoid_1 data[1];
 };
 
-constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(512u, 1u, 1u);
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(256u, 1u, 1u);
 
 static inline __attribute__((always_inline))
 DrawMonoid combine_tag_monoid(thread const DrawMonoid& a, thread const DrawMonoid& b)
@@ -80,7 +80,7 @@ DrawMonoid tag_monoid_identity()
 
 kernel void main0(device DataBuf& _57 [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
 {
-    threadgroup DrawMonoid sh_scratch[512];
+    threadgroup DrawMonoid sh_scratch[256];
     uint ix = gl_GlobalInvocationID.x * 8u;
     spvUnsafeArray<DrawMonoid, 8> local;
     local[0].path_ix = _57.data[ix].path_ix;
@@ -96,7 +96,7 @@ kernel void main0(device DataBuf& _57 [[buffer(0)]], uint3 gl_GlobalInvocationID
     }
     DrawMonoid agg = local[7];
     sh_scratch[gl_LocalInvocationID.x] = agg;
-    for (uint i_1 = 0u; i_1 < 9u; i_1++)
+    for (uint i_1 = 0u; i_1 < 8u; i_1++)
     {
         threadgroup_barrier(mem_flags::mem_threadgroup);
         if (gl_LocalInvocationID.x >= (1u << i_1))
@@ -120,9 +120,9 @@ kernel void main0(device DataBuf& _57 [[buffer(0)]], uint3 gl_GlobalInvocationID
         DrawMonoid param_4 = row;
         DrawMonoid param_5 = local[i_2];
         DrawMonoid m = combine_tag_monoid(param_4, param_5);
-        uint _178 = ix + i_2;
-        _57.data[_178].path_ix = m.path_ix;
-        _57.data[_178].clip_ix = m.clip_ix;
+        uint _177 = ix + i_2;
+        _57.data[_177].path_ix = m.path_ix;
+        _57.data[_177].clip_ix = m.clip_ix;
     }
 }
 
diff --git a/piet-gpu/shader/gen/draw_root.spv b/piet-gpu/shader/gen/draw_root.spv
index acecee3679ae274a8ebf82ee416e177df63eb992..1c11414ca510052be760fd68802677d698754e93 100644
GIT binary patch
delta 1698
zcmZ{i$xD_|6vprCmqUqYv?z!$$W26`&_<*{FbG<;jW!071>w+O^Jg|^lYPyTrj$8k
zn%ShApk<4;?fMtmG_Csm^u6)*K`%U<=RD_}`<`>|{ZunoQ;{iJQyM}NmV~k}*H^J4
zNv8JfENbZvA++spZRx}ohehE$w!OKtd2i}p{KwCP9DZk87rB|x58pkXmkGnzSTh38
zKVvLcml;lGr@v%ACrOeUE#AJWi<Ba;j$f8w4}r2Igw^;nxw?{N$!u<O$xyZ%owh>*
zdO6ni@_pOO_iZowE8w;med~t{d|S<rkEMPi1FjEYdROVUO!gq+0Z<Em(};ctaA!Xm
zL;e-rp@hSv2ZY0LceE62B*$To0QKk}g|~r^{M7g{&;-;cLUA^1g&YSK8vMcXvsgbH
zGIj!}*^oRoa1vhFKo6RFY~U0;zk$>EHlW@UvI||EnRh}k=;7Y>S<niMd6=s_2h=yw
zWOwSv*adtwV_v0xjP=5s)0o$)U(Y^Do$3;(1Qr?e0F%{!F>#z+`@gHWOssGdSJ2dB
z6IbD$8wb1w&pM%r{#3yvf4UFry0sqQ-Sq)i*`40y^;E0IT5n9v`gW}KzN5mX>(g!C
z@KzU-V6m?o;qTqvgcnY93r&5T27EWSftnA{cXJo(=esd>2dLSIJieQI@P)J4X`Jpp
zqJf@tx(9HzCJ<Nl5bnl|#kD?ytH&Yq-9Y@Y)>uE*MvV<%ZEBzj@||u6u|pu_I;vJ?
zpFm8XOy_(GSE~iirvDgN<uUR&#xuCJvAt2ayS0T8tiI!vW5>Y;U`@<_?xh&8U<`|}
z@K&)Cpb~hHF>{8o>al?r@U?mVB;1DLtC@nEqaJ_y67D&Tyt4)${4`EISPm@MmR`y$
zxXFG_F@t>#%s1DD^v#PkZ{XG>AlAHvTT>0J(KpY_8QKDItzJl+>K#0Piud?x@c?6L
N`L{bb{h_Ki`2!-N&W8X1

delta 1714
zcmZ{l%Wuq46vyw(Xs0DrVnJe2!bU_wqKl>>Y0{mQl|*bRX;Km?^?pv(TfJ($pF>fq
zsx?$mJ4*|ecK!t$ua(b7f0v(0{E~0Z_k7Pe_nvd_&CJBY_Y3QC*#-F!GGS_{32*!A
z)@8Dz8`o#s4~7srw{^64VXMNFa00uzt*dQQ;`jdJ=fVa2uFjq0=0XX+0mvU>W6dBu
zy~a>+S#BURQT~+skjZ3=gH>yn>>#BUEN07U>@JW^`McpYnGog?A1ltUo|Ty>Hdgl+
zY?m=YGx{{F?WTR(P5ZVR{W<Vb%D3LOV_e_X)BaH6w=m?A5Xx)vUvh<=h}(ff{G<{6
zEx<GU&KUB~=nf?uCRrftfqP~%!3uI5elJju{yumoc*my3_k&iTJ{+nFVHM;6u+ZQ)
zE}$FhXG6vcpc~kbJT}k+uWaBDntE*DFg)GB5quj^?+Jy;u8!hZXwb#I{bQg57;`a?
z@;FdmNt0fmevF-js~Phu^<(T5ye)}&t@=%zN8V{pgL+_*K^HJt{Rb0wW1U?4zpFSy
ztnwtzqN&Fw&cR(b4%i1TIH8Gds$i0vZpFH()&;z~^T4C*PVTapXbo8Fjj36`8Ed`o
zsIuv%<TNjOv$Z5x?CVDObGMh^l@nb?Qy-!M-^~@E<^%NIT*LbLZj4<8YBnN|@8&vu
za&68(PIm**Ku<E=O}JVsh(~q{?ui+T$9fyC9*5BP1meb8V<oJO8oPtFsXOc+zS3<!
z_8#yv;Xc*`ZEWxXyaYy)=?CCyOMw&X-vzE^tvuE~f|n9Mk9`b0@q(TCL!sc5)38s#
za$r%MayZe=_lHFP_Y1~81@*vnjG6Ngs~#H|fiFt=qwtdVA0N{*1d|NJP0Mf>YT@}c
zW1oX2FdLY^nudLtFMzpz4l#y(3CuUwhV;#gHRJGdSm58t-=bG=iyDAM`s2lq)8{t6
YhLu2U_YGVvW*SpV-)?_-bVhIH7sOrAWB>pF

diff --git a/piet-gpu/shader/gen/pathseg.dxil b/piet-gpu/shader/gen/pathseg.dxil
index 0ca0d185edaf3331ffbd83e41a1ba49692a9b052..3c8131532b027bf002a5b5c7a02d729315e7be6a 100644
GIT binary patch
delta 5289
zcmai2dsGwWwx7%-LlP#C0D<9|0P+S+cm=H&C%gr;!A7384xuU{Vg#!A%nLC^5K!P?
zt2I80iq{}$vB$OvpvF@`>cLj5)?ldzTUU$Lw%B^>`38bL>#lXz{U`Z;d(Zs#_t@Wl
z%!~A?boZp(ghbsn>4)e4IHGizIxxO*^5u3K1VL3z?$tC}=Lhf5><y4e&&8r3mR{85
z13SR1MuQ2dZHJ^tR}4LaG2iLX@FPw6^^TOkhkZh>A%9BBU7Tt4gdh@PKvO6LQ5XW8
zTtI69@`OawH@+AI{j~oQYv@1O_xr%_cSYFmrooNsu8UYcU4=6s$eF&KYfpzDKkzzz
z0CG_fnq~#z+ma=9S!#AECKkJ-7s?Z^x4b^;vWA%NlFxGTOK5&ca9rfi(dIIOhBo{9
z2WI*vXkk8S5c(!)U<75mh`?C7#bpC=LY}rto;1dn74maKWYMq;j;`aWh0NX}5uy^M
zKo4H38q7G2sTIDc=XG4jKXac6CgP8>AT6nENpSZ%!ev*96pfIOiT*95O$xJjHrm<L
zhkD(Ej}z>zHnm%*!JNQ8Y6NN&yxnN(&0)V9u3<;<-i9)diJYaVDS8I%=Qy!dtN&9t
z4HRn1Sbq2mvjR1U-s3_idU{*91Y@(oA7uh!!zdiGP#RJfNGBMj2G`pfEVLN$7K<P*
zU3c8=S7CcEr?TG!5nOEF$_breuVk2qAdpm));s2?O}7%M7zWYJLigP=X2~NJqr>Dl
z0;A}6Q*btyU83=}81zI$Fex*K8yGMaYGHH}YAK`L5QY<4?;(<bq;(@69VEosoTlSr
z!Y2+TMlEfQwUAT8(JMUEg^0EXC2y<o^E5nGSX}}Eu?bh_fqBC+MCeeXVXGa~aX1<Z
z&4(V4D!d)K4e51-2G}qo{S`>@nbk1nn8w`L0<C9Rn@)|bc1L}K7jP8}UB*PBQELwP
zRhnzII=v8ekji&=3bRE$)r(x@kd95tXjmwiv*@(iz-FB`7})lw+sIZDm}1f#v>G<V
zvkC#g0Z;+tG>u#Un%St_vIjno3va8rNFg<Q5l*cKLKGA@hl3UN+M_1p-_&MhqS28Q
zeNNELdJRo2#OEyJIf%pxk~+UVOD7WJF$rBvL%XmPUUQ@4GIN|sPlLg8#gZ(i^W3)n
z;si+LlI0+pIq~<F8j>RDA!2ai?;J$xa3RFhYQA(1)7aB8sKX9<1pIQ^0W`ZZ9w$sX
z9b?gCoW--8MLg$ExfY6ZzzIkAYFQ&&ETaYQI>iM*ZlFN%uG192T{;pliM}sSZ4Or<
z<@6YbIkYhRcfSZNAKP)nwU=G<Y3`mvM&e2%qf>ii-kqZc82@WuvyzeJ3uVdl<(aZ%
zMi<>eFoa&EBqoQhn$fES!!QGerh(AAl(a1b^(Z^ySoXw~nzVE!$zaaNRhkr!@-fE<
zvD>twm}IbQFVuK!@^z}%I4#duk;bCg<X1hm7CKdSPRpC}fqcMe4#I~#gsS5lIpp;;
z^$x)xeWyXO`bAT$Ox(_*gjY;~7QtGCuX}&`NldAwlqy;~Co&zpLip~W-uJngnjdg+
zqiv&~)Rc@Q89l(%jMv8kHud$8_X-g|;1BO3na|M3Qi47N$&Y&Ggz-qm<s*F@X_eaS
z+cy;XcCD;`-Gnlc@F8ViRjpu*(fT-cTlMB}T^H9zeKxhM@ruDy1-O-FqfL0pZo|Rk
zFFa>tEfc71$$<iEr>nFIIDv`Lm#y=hb(kHE*9z$0Q4ZsVz4o~&>GkA?G#k?+yR8Gs
zjXpC>%Xit?m>voqZ*-Gxn_+U0RQ;}^+;fRbbE}Vy^p3vYZYBNc3~9j!+h*M>Xqmvk
zyL}C~*9y_&aA~y7XKSt3iuKZ+)4}NZM%Rf{bS~b({L^BOdWs1Uk&uX#KNDkyp`;XN
z*Z_0q+b7IH+_3C^{WwE67`d=X<>}&8ZSw_9tACK6C#`mxEV$hDcp62f=U)e8WfM9t
znl29wW`{ba(Ni!k2w{Y?Gd@}`OrfZ?_<Z}fr>W`n6tz<l5`)A?PnoQJ3cE({g`2F{
z2=;7v7)f-r_a-VQhc!{MsGh`cRStEpg_lB|=00?Y7*soO7KkQ{?X=@)gGU?^ql=4v
zV6t3+Q7*o)d<-yL9-f5#-S*gRX;=Z&NJMnG*>c5CTK>eci{eZ0cdJ_w&Yj~1)bH3w
z9pyisU0Q5V;Zq+?8?`iE9BezJQPu&KB~=y(PDz$}V0Ndsln4I8Gh7Ed-Tc(wb}Al+
z+gUs5>mtEX-)f${A-@K#KQGX{PjO1u&T_uXJ~)-i4WbTy*!VAO+ko*yY@5Zn23wbc
zTdrZ^#s+tVP!M&Mh!5mWiZv>AgD<=t<5`P_+%4C<q8l6L7YhXyH*IQlgA{0;tWgRj
z71RdLHP`yp;a(szNs+lHk23vHJCKovtM{A|_3O=442*tI9*T+Y8XYhD0nb{-za41z
zv~#_~c+uBhmTR#K8XH*j?U_Pr(IH6i1AVGaHAB`n_x#OSvhL42qI+#*KSqXe#5UW1
zA5KriBnJ8=5Vrb~7&Yh)p&$`i)Yx!DBd}Z>wmOa%!8Cz0bE?gNPKu?R*-0<?^hLVA
zZPLf)Szp<CY~K^zud*3>;WU>c&LHyw&?lYg*n=YY1IxAQg-s1j8c5)e(Be_MVy-}k
zMC%(^JN*RCe<sDBMB(k39QxFNQ|@ZJGOnNtkpSg!piEB&kArP<pikl6lhW;V4pYMT
z!1mhhnYsMuZ35<`coEkcl|gbUu}N>yMwcvVVEg}A)JaQTj(F?lW|k6J`%^u9_Am1A
zQYn+(yvzQ1Fp|t6KT_I16&Dy8&vS*QH8XkMr!(Y$sZ#1TUUaF^{%I$nzA=-J7OvhO
zsd>Ju$8rr$j1R<~iqrHB(uD1GYHJalemD;<H7JO=(MB{B9Zx{k{V!mc`}KLC`Mx(|
zU6V?8shd+4UM_O`=z4v!!^5YsA9q^JE3B&}H5_JwJ4lI&G%87Q8>ytPErKr2Z_Btl
zdDTecSq7op4SIWKjlZJkt(u=;VR-QSo^qw6U`!k<lirxnOA+8T%5A#XDW{G>W9+)F
z*q6(`l}cph8xxlhV}OBEhexY`!_j_|l=jUtf6d(RbL@tw$ln}HuF_bMVwm#F6%#%6
zpV2q^_b<8IzvTKsFQ@>$YqBvPLnK1yfqIdoB1#feQ2~|uu0{M&4scsj%ig>COqnF%
z-pG=B#|OTk8Bg?3)+l3}7gQ!jn(r|;S4cK@MXGY_Os>w7S!>_luD?#-;y=FRoAD)|
zAM}O_h6ZM>ozY2xogdd5CEL0r^C~t&rP8WdYZvx@)1EmjNx1)P$^HJ!<khqCOGDpu
zWLhOJ$0R$yl|;U=Z*Z3DT%RP|`NiO`5|!#_NmPXbD$SD2zI0{p{hmzjf`p0cr4u6q
z^)%zEPiChsdCh=CnJ1XXq81#hh`Jz9A~6D^jz7FNpAd+G|ImAPST~K71u>;OO?IcB
zOapb%@A%ohH|p)2H%>%duJGIAReW-@?CmCGaB68S=E?29p~6f?n)t;41W@k{3h)4M
z(N8#JbK^fAG_UU6oeYX20#`eg&});vM8kpwXALQ)fmEiH(RqyD&up$WIA$B;^JxvH
za80a?HpNR>1xgWja%VHhx;v}gM+t^C6>s6lv7?&jSs-ic(TAy{#1)}On_!g+r0wFF
zM^CDs04;5jcXE|$;BSa-tcKWWB3lbHw8;(|b&ICDEWTw=gH*2b`|JOlxZ$_EXY}y@
z<UQm{Ui2-JW4hWxnHUSL0kz?3R*$AuI=ajk15UmwW0weT$D1Sgy`(`QyC`SsG;Ct@
zElLUD*hN(iQ<v;;M3-Rp+daJ4x3DR3qJ7wC@vY3(=RV{v2$#>WUJ<q9DBl#kx~8(n
zx6N0C?+xDjAk24ojERJN>I5HJ#njDh@2lz77VY+gy<wfbx^2~Dd|@2?&*}}-&$()H
z9;9fsdbUIK*B{k*YDKUvuSNvvV#$6hRnU9HXnjTPDVPJueulhbIIJtNDPn8?QmplN
z2Z~l9w*R1ym#WNAM3@MNc)7I}`*I>QSa+7F?}p2EbC^*Cb!UThMWC={MvA7Zb{$pa
zGQ~{PG-RzOMR&RFJbWNhPjl-ar~Kpc4=V(*`}Q4!T*cKly%>FHk`*I9x#78ZGz;2;
zr(NHx1?7<%HH@e3J^naYos50h9FuAW-jasMet9Ex*=84$AG-X(F<_S2PK64OVrVt%
zD%f)+BG%i*fHSyDU=O66{aFHM6-?NKP+BBcS0&Tg%JbZJ>%VMI3+WlFdC9Vntx}Ff
z7|eB2FcKpUER$)27av&$CQ}yD0<&wfYC;xYpx~xc)37#l@g)lG3~8mbd&3r+DR}i1
z1*JSus3pL=j9JdYYWy#%{(CBvTN$VQx%M0^UE_Y!fN8(4TKofZQ;#ncB<g|md1zD|
z0#2_NoZit>uTGEh88g-#iI*{lSfFdU330SM9o2yqTd%iXj={xoRI3*FLD|{@olP<R
zsMR>DSp9+WRK6}xtBS@aaBlSEt|%cRa&|b0yc(MR8a5)4?E{Cp;kes(9zF`y2c7l3
zod;99c|2OZTl|0jL(S9rnQm*<@vDmuf?NDXFy#qChHARukxh3$lhdsbb2;#fN%An|
z7YChzUv!vp`b>*!Y?jY&`lqeEc{#FlOJ&B7(V$x4Pdc|gn%dvYHH9a78Z|$<2!3(u
zAoU?l?DWZ?kbP{D*O%EMC(NRMa6?WtQ*g-%3Z{!gR3R-CUV33#ofQ(&O2OB!P%s$2
z8ZPz3<d~PIVU-aD;o%Be;@)8<u0Wl%0(6~7CJrvOwwgUaU&8h@tu>-u^fHl*W|4Z1
z*Z2TorAN24Y{Kkkmw)FJ*@Q$qX*OCkty@cS$BB&i<nRi^iNmT;>NfWo6gUVRY*Jp7
zz2GjljgMogE~!sixD(21jmt`wW+-yN>`8y;cuThOu`+Qun^e-cU6_(<*Q-pr$huGc
ztYn$rg8IXyUcfNy6}$m#-!LC+-+;yoY~OSR_3T#eE@{k6rBlDMX*n+USIuZ%@bg}2
zIK?eM(?zck4U2{4Lc5IycG>?wyemStFwj2E^X@lDna1D-WgRMhq)Zg%gWi?2tK^P&
zG=90<gceBg$ANY`A|qlAB<=<ksFce59NL>^TOK`oA}5sp?ACD3j$1k63xnpv*<lXB
zPztY#VaulG?3EYfP(EC&aV`Er9?)^3k7Hh^0X>iRh*UJw6yI2X{njeY67w2c)VlY3
z*a7xj=%ni~xFIOGA=a`z-Nm7C5o6s;e^7OADzVi!-+X*#P<u#w!&a<)mYgr)uYfY8
zs#)7lvP-^`oATjxov!ONw`|qV*z)|t43cr!cE>Ck7j}mBKy9a2*r#9~AcCsx=Y93`
zl<;Ep?b@@qL(j}am#Tq`3s+6r^V%D;rQxGr^53Sctvx+!?Z8*1^~=cu?fTKy^{1Af
z?X5D0YrL^aD8bv>?`5S^!w%08TG)W*2egv9D{kPhC(PdTKHmCCv{&Nw1-e1FJ%Ik&
z=W4qx+&UZw6^Ge><=Fr`*M9zXTX;XEns=|{j1?x+*Pc70{)p@+X&?i%|FJKkM|+OE
ze081eBT3L_#rBTjrm!hniQBhwEIG%_FpZcCL<%ES9m!Q=t#6LzG`{sHz-H^`mEYd7
zbDer+izj~eJ!R^2;XGWNOSNq!{o9N7oZA{!b?!H(l<5yf){xF;P7M1r`*@qn8rrti
zXDex<m=kVwcsNxwwJ62|^1tPe<m_N|(YMI%gN9c_UA6=twbmPd?I4cTc)~mc<h1`%
O>ZlJY(EiK-(Z2y6c1tAy

delta 5314
zcmZ`-d0Z3MzMsi5Sx9hz1j3#WQC6J*8bmKn009AQ)PRUq2M|T27^zwpYBs=#pny@S
zS{ua$t+r86v88<pU_nc1RO$uAHu%&_t!;5>FIKO;XR_4$-sio4lKGzB%=h;@=lsrZ
zIb#{GGJI2VlalWa7uNKx-Vk=?!{*m#WY{M#1VQEXz7;U#+~Ezdn-IqIC?^(TF(oZQ
z6n6>>Jz}O-H$xhF%M@BBecEx2tdW0x#@{M`&3)*;ogqxgot<^xAA$&o4!uAih{O=!
zL;-Ct7*9$ze=p)d(CD7atiJzb-xEYR*b?ow9j2g~mW!M`n#z|BL0+_S)Qtu~V(>Yw
z7vdG*8bbjt+B;9in>(AAr9pTZ`HG~gd*2%5t;DDC@>rfBN&UEN?^^svgh$}MF$EpG
zknv%17-4&c@WL>{h{IIyY)E@Cj<*&+u1GIZqzv=r`TX2)!=LahF%*h-3Sg#II75ZR
zLWMNp>p%Y~sPRw>nH@g}EO{Sa_@fq>0t`Y}kd{z>p!E$nDr9fQ3hE#s6Zw1iA(<yz
zWn?(TM<{+_$K~v;PVvHsW0v{sCNmHt;On}sBtc_`UIVLzg7N22<Rs=LL(FmPSdZmU
zR+H*~Aj~WOjmQba2otavuTsLIicxkG3zlDtT*87CYDPkVnl$Hjw4uYRv7<N%QrM8y
zmcSmPR%-R2QXv_ZUalFZjj-Wq91cXY2w}!d<|92dP=Y}yO_vgYi<(s_?(4hE5RP&@
z%g{=r_cAm4H6Be6^1jIg{ixB~rr}gG`?VB=Z>8i3%|^=GxEoENurWT<ipJj_=Hy@%
z71NCtc2(V`?Qyvl)N|B`L<*8Zf^HE5>aY{AV5(5R2CCxfMgHo1YEm--hM6^RV<;1e
zfTlqARD_@z8iDk>d?TffDs->Xu+?03yZ}hq_Ce)1FzA>~ZW1M+>&R14+i=x1W)HD1
zAJrbEigK{_u0&d-V5v5^g3p<JeOxEIYX2ysfZmz;Dp^LynjwS_>=r1XCN}eQosrEx
zZ8BD~R<KA_sLCb;2B5IJ79bBh5ZVDw0>?tpC;lP|uHrH!sNH&9K9~;^6Ov%8nw()f
z0daIJ0Y^&v143FjDQJyA05P?)8(#NNFNYVH9#<reb}EyFUIg|v!IQE65c2D&H3F3o
zyr`-2njmts%&wV0!#y4Xue0y?&nX%GNk&g{#~BG`Jxl=(RGet>Y(rbSz-Akcx?>Ip
z$jfc-nnc4S2T<rc1Fc42r%d&Z8FsT|0ZMa}Mk0rqm8VcDUp?Q*1|YriGzoG{Y!yIy
z*>FUcQWtNe3wmROY4ZOjJWX25<o=Lci_K?^=nOIHc2SUnjm=t&MhaeaM6)B*|FN#w
zaq=tF>=gMGTDAhqRC5$ihXwfxrxO7s1OV6Uz~%rq0B}y^$P@2BM!l6P(rW$~JxA@D
z;c<!2Yn$hK`?wy-sdJ5=TIUlY^Vvj>U+Ahdpsw+==X@s2^C`E#f0i7NlEuhjBm6Y&
zHcmH-RL8gq$lDIoNKf!8DqnX%NDgvJw6E%?s^pS37Ad^~O|(m)JHjUN6W$H-dz<L}
zW7%c?Qs+XiJ(D_zc7{}SFkq=xTH|&raH>$74_%yQu2?JVL{e-VJoNK*7d+t{tAD}}
zs%HoSVjf2&pJ2{1SGh}R4`5_NeRsAn2pdu!+YVMcUgainaQkf?K^Ad7C%Jt%)14=K
z>+`P`@xFA7Exuann({M`pfw))Zr}t<u#zd4VW-2~P21c)iwFV^qqLY>I9bmX>w}yw
za(mmc?j^KGPNuz^Tugtet_Qv6(dM9ftkhi!2JUH1k*)gCMOv>fADx8-J4hdV!Em})
zRmktKpXGhhz382No+(Zf&ZcjxsS=Ekd<tXmO3Sw8Hm#+4WOqAZ@}SZ{M6xcYr6{P}
zq~&oUD1uqvqQX;Rlon!EUHVdNAga^rvdp5<yR1vlhMe~})<crjg{;(6+sTM}KsH>J
zX!XDvi_}MrPmL)=7X*U{)+*Uif0w!3PF+{uLZsT6`|ZrKIC-osSvQW87_|`^v=lP>
z@Ne*O+`p-?=O!zh!Wb7APo1Kh9`}p6t+N9`{b!WeIQ^UvtC6a6-|2_19&Ogd%47A(
zV)vRQQ3Ry1h<yNq*^N6HD?7KT?w>uCKx|DC`v9@$c%R)%NV>?lGE2`DimdEOPvwG3
zxHu2q?V8&9kO{wf)LboAB|*G;v1^w;H0&s+t%w0Wu%VIVT<H@Ni{}T+4v?`SoS-<@
z`M0YFY-#2RWNeu8fs504Z^BED_N!pIc*iEVD4V->sAy#*C4;-Q&(zJ?dd;+*v-K;}
ziV976FRG%xXcQq&C54>(g%mj0i#a9;F?XYI;=@BgECW&<kWS?OB^__>7Q%^QgaMM0
zRlW06Scl$1rnb1@vWRM|KuKn#+O|XN<QhQ$83SV{HFuz$+?UqK0_<yHw&tBOVep-Y
zc;dxxH032tRMm?{+}us0B;I0;o|~)w)=BnrOk}N#>~m{|*h%(mQtJ)WsoR`+Hm*xa
z%`muubCjhZimi;NlDJXo3xVI5rkulljVDe(M<p8dZq{lj*sC9M89n0G7UwehrOI=X
zOA+_KDl%PG-<*8!2HHZ~mtBAcf7~TZoM78gp4H>^Lg33BmBgU7LZupcB8kU!V~w_n
zPbI@|-7j1%G<Wyo>TjboMH_8EI9&;Zc*Il55ykyNB@o(h_3rC%a98(;dfdH2?dhKG
zitPI6-7hpy@G~l$*h0I7cSGa6(mcL<^KW$X@|P>7SS2GJT?j9G&tyMrPnI8BmDYdd
zEcQbbbt8SF+a+Oe3w?3jZaSqr)t={Vz6TpvSw8D!iTm6x3xnNgV~%`ed3S2KSK~7M
z&hjYieqo0&SXVJ77~j>+WG1HKqaTEtkqG2|5_OHS0G||R`rqKu`Dh?9P8vGzJx$fP
z^N#JS5XQTk4-VztCnhPI{C5WQhL&zD98e1lWJ;`BP9{m%7bN|Vc!W5pyWGQF%kf7o
zLmc&QRc&e_K4j{{qOnO(2c_k*jJ~ObAGDmxAFJ6Vhyz{&cS~34Cc}3)C7n}Co=u(-
zv;kzR+QZYK#Ywl#dKpy(98PJxIrWvNxfYTtsZ#N~zd|J*MN~1u1%7GX(RXJoQz+FA
z_0$d-dauH!qAqTYgE=i>9W@5N%Uqu<UH@*3lwtNhIu$Fhl5TtG$KKi=+D7p}?cISn
z*8&2e#npF6ZKSg11=8_ewpz7xlTI4>u4)5RA}MoeJJol$G0P~`em_$CeQWO^Y<ghj
zIJI>rY}UoVKQPNXrRC3J-UHeUQkS-$9@k!@Z4-~x{xCMDS5e`zV(6#QZ?4f)&w^FY
zq$+W2=v@~VWBU7@G4M!<k(!_~#>P%W%-)r;F0#k3YQLr(73+76=;wYjB!Cus54lWl
z^-93J>UQA?$C4AI3toCi<<C5(=gNL=ray;u=f@3N*vd5Sds>v|QOu8|@R?YqF}lp?
z`@N)omV4aI=@ZZQXjYUnW_eGu_WpXT#{|n>^{ODSm7U*m%qCd$s+R<oZ#lT8=NE#0
zujx)eNoslSa!}`0z+p#Hjp-Hdm?+P#vscp|msg+)bb6kw20xYl2SI<_#2;n$dKQDt
z4bE3gmBWqj!?Zz=b<-N_K?>ehF>esRBJ^tvE&^G*nAm0jS^KEw#XK@=>k<2_1Y1B#
zq2T<y75Unkn4pyTkil8udvSd3eN_5l&&-_x_fl2QSJo#FJ^uD}RGHrY_3%eU?(1&O
z{8#c7tXY8YJB<aPHZ0ty_HPy#qWK+!u|R%N!PIH0WrOoY9B>SZGAl7qBpQAug_JrA
zu~k)x0s4e+UQ`(~knia|t)h?3-5HRyy;5J~e&l8n-NDMyRp-m$nu6l=?fw)$if&kc
zFYHTCO5iT6?p~zGx~K{p(I%9pN5*e#6Ag&4Hb}n+`PHjRjLjLd|70+KrwCWDzSOYs
zQA25)2o-~r$E0G+nW-FiP|zC}H3=xX)m=8Se)pT&Ai|oRf>mPQHv;JaAXvcy#uwsd
z+&xiY+fy49p!%doUHgePW`KPwJ<_&uK#ft!GpmmQ#}%AnY9R?EVI)9x1oW6=V<Meg
zi=W04r4uMEY(!q&^jrAvtHh_Ku?Y`d*!6i>2~O*S`d^ys`vd7F{Sm`2VIUGvilK~c
z57Qs*qU3g|Gq&)6vuNU=?3t>K<15qPsDj+%iEN-r?FT}|Q4K=uBgi$EqPk6Tum=`y
zmoTUw?1_>p(KDt9ngv5R=bnrU`hz&fD{S+4hsaT^AKyFBQ+wK=^)DU!NDUmJU(L}{
zM4(F;xDiedt!Yq10Mp?8ig1;Y5ELD1Fhl~Z*}o(lz~g5c`bh{MT28_T?G1`3QhKO_
zggm&3^7f74Dty10h%fl5`csN*yEy*Su?TeiJ`wktu=$hyt-FZr{#qqP=caBf!gR4{
zl~|Ys_VPf+m~$`tfCr(yc(`s@RRw)$Ccu|g&}YPx4cUalrB@g0JalB|R(zke|Kb#<
zV8vl2>o37b`2N#V8bi}{nq9Drx-6bLU-;MK%r}oS-aJkd>aKl1J5sGXTYLBu%H|t>
z+7CnJIu&s64C*EU>FEFWKO3J-%AL(u`szd1Rt>1*D+|cTKHuK{Fe~l7&`f({B(DT*
z(6})c@Dj~5;3b`YH8QM#AS3I}*?VtJt93Lc$oyXnX_ueRRguMuY-aukeKTjX;NF;4
z*ned1lk1a#XMNyz%2=Ulh&Tm#ORR?3bRMUvroLf532lZ95*qe2<dJaHkV(QrdmC1h
zP`x(`KwSvY>)C3-ub9uldB_?}F3Bx1?z8xD6dcswRAfRxH0eN8rM_n@INfN{90KoT
z+Q6tfpjB2n0tShbGXMKkK~ZJK(8FvH_~yUG+8I!M^<>p7_txhNx?&i4l#jLGY)d#1
zRWcAppwuZVkAuUpNMli^tm7?XY24EU=}-2eBT=0J?K1iHyjApWQ2&?G5g$nMHN;pa
z<fzh?9~|Y9KYP)~biCKI3n&(vHBZ6%$=)W{FQ5mkpQE5!R-Y|UI18Uw*^4S=Qyc<r
z@!mIrihZOY6Jt;VsUeIqlyMqN>i>OIIp65y78(~4*ihYD6X*z^xdleQtKA1|ZK8u#
zb|cP^ot3S6n_gn#S6E8RtS4_4P}x=MU=y9uyn%1!JM-K>^0pqIuUW?3_F&;Q{(?IJ
zuXbFxr=}3R`i0H}wWQ=lb|(EuXhlwOrT3id8+i$;A=K;*O1A-7AkEPukJ@f6dE0w{
z2g(?GbQ_GeHk`6ih}EoUfutCm5UTs7jGZSa{ho5S3-UE5oz>2d(e1aR2Esv+V=p+V
zOOJ0@V4<6Nv^<yYwWo6C($9vr=B&Bkom*?~%`IZp(i6Wef^tTE9f$8kd(KZMwQlND
zk1h~hMzWoVf3hX|Blhyly6C}%v*-OB;pHFgIcqpkVUaFvO-oI)a)36&MVGdr&DU?`
z(l3P$-FoNlzP#@u>U?@V2P`;4*bmR`h38tqqwO<Uph?vcU@pZOom_Ye8B&9WWLUs;
zWUa8g)8m~{q+?+cXyG>1F)WUoqZ!X0<dqS!C6<-WX#eu|*^BH+&<Wl@KVxrkOXYZj
ze4y_5bWr>87`ta@yYtFq)ANrnj+!a$T32EHuh6CI`AeXzz%NvEG2#Y_CH%$0(!uZE
zr$0Kr)OkxaLzueseL}O$rR#u4=BHnhx(@uV_0Z8lbTXztzBG325@k#~Ip1-66y&vF
z@;sBe)fVe;%sJIx{Jog&11;FI#iF;6R~Pb#oTW!+E!Dg~xF5@`WBewn7xa3PhfR1N
d4ja&%qgGppHW9dTX)GF`O3C|`J?Wv?zX3^+at8na

diff --git a/piet-gpu/shader/gen/pathseg.hlsl b/piet-gpu/shader/gen/pathseg.hlsl
index a9cee25..f7c9e2d 100644
--- a/piet-gpu/shader/gen/pathseg.hlsl
+++ b/piet-gpu/shader/gen/pathseg.hlsl
@@ -72,7 +72,7 @@ struct Config
     uint pathseg_offset;
 };
 
-static const uint3 gl_WorkGroupSize = uint3(512u, 1u, 1u);
+static const uint3 gl_WorkGroupSize = uint3(256u, 1u, 1u);
 
 static const TagMonoid _135 = { 0u, 0u, 0u, 0u, 0u };
 static const Monoid _567 = { 0.0f.xxxx, 0u };
@@ -92,8 +92,8 @@ struct SPIRV_Cross_Input
     uint3 gl_GlobalInvocationID : SV_DispatchThreadID;
 };
 
-groupshared TagMonoid sh_tag[512];
-groupshared Monoid sh_scratch[512];
+groupshared TagMonoid sh_tag[256];
+groupshared Monoid sh_scratch[256];
 
 TagMonoid reduce_tag(uint tag_word)
 {
@@ -360,7 +360,7 @@ void comp_main()
     uint param = tag_word;
     TagMonoid local_tm = reduce_tag(param);
     sh_tag[gl_LocalInvocationID.x] = local_tm;
-    for (uint i = 0u; i < 9u; i++)
+    for (uint i = 0u; i < 8u; i++)
     {
         GroupMemoryBarrierWithGroupSync();
         if (gl_LocalInvocationID.x >= (1u << i))
@@ -547,7 +547,7 @@ void comp_main()
         local[i_2] = agg;
     }
     sh_scratch[gl_LocalInvocationID.x] = agg;
-    for (uint i_3 = 0u; i_3 < 9u; i_3++)
+    for (uint i_3 = 0u; i_3 < 8u; i_3++)
     {
         GroupMemoryBarrierWithGroupSync();
         if (gl_LocalInvocationID.x >= (1u << i_3))
@@ -575,16 +575,16 @@ void comp_main()
         Monoid m = combine_monoid(param_23, param_24);
         bool do_atomic = false;
         bool _1263 = i_4 == 3u;
-        bool _1270;
+        bool _1269;
         if (_1263)
         {
-            _1270 = gl_LocalInvocationID.x == 511u;
+            _1269 = gl_LocalInvocationID.x == 255u;
         }
         else
         {
-            _1270 = _1263;
+            _1269 = _1263;
         }
-        if (_1270)
+        if (_1269)
         {
             do_atomic = true;
         }
@@ -612,37 +612,37 @@ void comp_main()
         }
         if (do_atomic)
         {
-            bool _1335 = m.bbox.z > m.bbox.x;
-            bool _1344;
-            if (!_1335)
+            bool _1334 = m.bbox.z > m.bbox.x;
+            bool _1343;
+            if (!_1334)
             {
-                _1344 = m.bbox.w > m.bbox.y;
+                _1343 = m.bbox.w > m.bbox.y;
             }
             else
             {
-                _1344 = _1335;
+                _1343 = _1334;
             }
-            if (_1344)
+            if (_1343)
             {
                 float param_29 = m.bbox.x;
-                uint _1353;
-                _111.InterlockedMin(bbox_out_ix * 4 + 8, round_down(param_29), _1353);
+                uint _1352;
+                _111.InterlockedMin(bbox_out_ix * 4 + 8, round_down(param_29), _1352);
                 float param_30 = m.bbox.y;
-                uint _1361;
-                _111.InterlockedMin((bbox_out_ix + 1u) * 4 + 8, round_down(param_30), _1361);
+                uint _1360;
+                _111.InterlockedMin((bbox_out_ix + 1u) * 4 + 8, round_down(param_30), _1360);
                 float param_31 = m.bbox.z;
-                uint _1369;
-                _111.InterlockedMax((bbox_out_ix + 2u) * 4 + 8, round_up(param_31), _1369);
+                uint _1368;
+                _111.InterlockedMax((bbox_out_ix + 2u) * 4 + 8, round_up(param_31), _1368);
                 float param_32 = m.bbox.w;
-                uint _1377;
-                _111.InterlockedMax((bbox_out_ix + 3u) * 4 + 8, round_up(param_32), _1377);
+                uint _1376;
+                _111.InterlockedMax((bbox_out_ix + 3u) * 4 + 8, round_up(param_32), _1376);
             }
             bbox_out_ix += 6u;
         }
     }
 }
 
-[numthreads(512, 1, 1)]
+[numthreads(256, 1, 1)]
 void main(SPIRV_Cross_Input stage_input)
 {
     gl_WorkGroupID = stage_input.gl_WorkGroupID;
diff --git a/piet-gpu/shader/gen/pathseg.msl b/piet-gpu/shader/gen/pathseg.msl
index 0f60d4d..9708585 100644
--- a/piet-gpu/shader/gen/pathseg.msl
+++ b/piet-gpu/shader/gen/pathseg.msl
@@ -156,7 +156,7 @@ struct ParentBuf
     TagMonoid_1 parent[1];
 };
 
-constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(512u, 1u, 1u);
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(256u, 1u, 1u);
 
 static inline __attribute__((always_inline))
 TagMonoid reduce_tag(thread const uint& tag_word)
@@ -432,14 +432,14 @@ uint round_up(thread const float& x)
 
 kernel void main0(device Memory& v_111 [[buffer(0)]], const device ConfigBuf& _639 [[buffer(1)]], const device SceneBuf& v_574 [[buffer(2)]], const device ParentBuf& _709 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
 {
-    threadgroup TagMonoid sh_tag[512];
-    threadgroup Monoid sh_scratch[512];
+    threadgroup TagMonoid sh_tag[256];
+    threadgroup Monoid sh_scratch[256];
     uint ix = gl_GlobalInvocationID.x * 4u;
     uint tag_word = v_574.scene[(_639.conf.pathtag_offset >> uint(2)) + (ix >> uint(2))];
     uint param = tag_word;
     TagMonoid local_tm = reduce_tag(param);
     sh_tag[gl_LocalInvocationID.x] = local_tm;
-    for (uint i = 0u; i < 9u; i++)
+    for (uint i = 0u; i < 8u; i++)
     {
         threadgroup_barrier(mem_flags::mem_threadgroup);
         if (gl_LocalInvocationID.x >= (1u << i))
@@ -615,7 +615,7 @@ kernel void main0(device Memory& v_111 [[buffer(0)]], const device ConfigBuf& _6
         local[i_2] = agg;
     }
     sh_scratch[gl_LocalInvocationID.x] = agg;
-    for (uint i_3 = 0u; i_3 < 9u; i_3++)
+    for (uint i_3 = 0u; i_3 < 8u; i_3++)
     {
         threadgroup_barrier(mem_flags::mem_threadgroup);
         if (gl_LocalInvocationID.x >= (1u << i_3))
@@ -643,16 +643,16 @@ kernel void main0(device Memory& v_111 [[buffer(0)]], const device ConfigBuf& _6
         Monoid m = combine_monoid(param_23, param_24);
         bool do_atomic = false;
         bool _1263 = i_4 == 3u;
-        bool _1270;
+        bool _1269;
         if (_1263)
         {
-            _1270 = gl_LocalInvocationID.x == 511u;
+            _1269 = gl_LocalInvocationID.x == 255u;
         }
         else
         {
-            _1270 = _1263;
+            _1269 = _1263;
         }
-        if (_1270)
+        if (_1269)
         {
             do_atomic = true;
         }
@@ -680,26 +680,26 @@ kernel void main0(device Memory& v_111 [[buffer(0)]], const device ConfigBuf& _6
         }
         if (do_atomic)
         {
-            bool _1335 = m.bbox.z > m.bbox.x;
-            bool _1344;
-            if (!_1335)
+            bool _1334 = m.bbox.z > m.bbox.x;
+            bool _1343;
+            if (!_1334)
             {
-                _1344 = m.bbox.w > m.bbox.y;
+                _1343 = m.bbox.w > m.bbox.y;
             }
             else
             {
-                _1344 = _1335;
+                _1343 = _1334;
             }
-            if (_1344)
+            if (_1343)
             {
                 float param_29 = m.bbox.x;
-                uint _1353 = atomic_fetch_min_explicit((device atomic_uint*)&v_111.memory[bbox_out_ix], round_down(param_29), memory_order_relaxed);
+                uint _1352 = atomic_fetch_min_explicit((device atomic_uint*)&v_111.memory[bbox_out_ix], round_down(param_29), memory_order_relaxed);
                 float param_30 = m.bbox.y;
-                uint _1361 = atomic_fetch_min_explicit((device atomic_uint*)&v_111.memory[bbox_out_ix + 1u], round_down(param_30), memory_order_relaxed);
+                uint _1360 = atomic_fetch_min_explicit((device atomic_uint*)&v_111.memory[bbox_out_ix + 1u], round_down(param_30), memory_order_relaxed);
                 float param_31 = m.bbox.z;
-                uint _1369 = atomic_fetch_max_explicit((device atomic_uint*)&v_111.memory[bbox_out_ix + 2u], round_up(param_31), memory_order_relaxed);
+                uint _1368 = atomic_fetch_max_explicit((device atomic_uint*)&v_111.memory[bbox_out_ix + 2u], round_up(param_31), memory_order_relaxed);
                 float param_32 = m.bbox.w;
-                uint _1377 = atomic_fetch_max_explicit((device atomic_uint*)&v_111.memory[bbox_out_ix + 3u], round_up(param_32), memory_order_relaxed);
+                uint _1376 = atomic_fetch_max_explicit((device atomic_uint*)&v_111.memory[bbox_out_ix + 3u], round_up(param_32), memory_order_relaxed);
             }
             bbox_out_ix += 6u;
         }
diff --git a/piet-gpu/shader/gen/pathseg.spv b/piet-gpu/shader/gen/pathseg.spv
index fc63eb5375b579bd2fcdb56a795572c245b9a9e6..37c98470f4f2f05b53b865e6e1eff5a7d53a8792 100644
GIT binary patch
delta 2552
zcmZ9O%X3sk6vlfV2_zVBOl0sCK$IB47vh6gG$KTlh{_``iHQn3MU>LAisJeoP;pi+
z+*``Vnw_O97pz^YMOiKt5tXNk{(e()DqVA{PWAVFefsp*J>B=t*Qd*8PM7;C<#kI^
zDy81EG#y>&n=6&ReEX07qQ5Jpcjn)ne=DWbm8$9X4!*X7-`ByP?cjSm_<;`oYR1P>
zwRvvflS=7G^TSmqx=W?zZ-cwbJsVSZdas<)k>>c?oz<2@oTqKLo21oY&UKh`9p+qz
zIoDy%b(nJ<%17Cc_sgldV*Pie(n9m6p|K&F;eE%trv9p?^e6EPY1dLYT?Aikx0e{)
zi+&010`VhZ{_n(<R7qE8eXg?{b1thBj=$Wvc%Z)sW~+%Fvaf=FCO!u5U}QJ#HAfL=
z$&p_NtK9&9fu@VU;p(Tr&i*gh$Jy=shfs5NF|~_Uu^W)H#7FJo|AEz3mKF`rmA*&Q
zq`yoFUatUXHMs#;2CJWJo34P>Mw#UV{Z;TW+T(<CjHIH{5>1Sic7by<-Eg(R%*+ya
zm~qFJ!q;F%J?iv;)noEra7=E^5q7BeW`K#Y^~)Lyhx&_~P?iy~*?nNOZY+Gxg5B)p
zaP`=1EBUzD_N^dlgqtmPx?frI9625%^X)SyEZ+iV>Afw_Zt_5*G+i%lMIIn5;T4u|
z!G`BxC3uEbJyu!=$4V{fI5l%)s#V~8mv4isd7iA}?03L9hmCJ1YQ>#kxwX8M^6+pM
zSP$yqVGtZ1tnD~8bHl@GaPDCZT+M9`4|l`$uz#{vtOLuf<%>Wb9_|6p;z2z;+zSp5
z)^?nlx#3|wIQK9FS6iNW7>4WNgXvmvKUi)pUq15iumP+G_3*F}93HIgI5l&_!vo;l
z!-H_OlXx)SJ~i;s))_jFke?yXk~}WAt;Id357Fnj^M&vT*vB)f{xG5DITXjHZ)%KB
z)@zTVY|7Nl9>%H1P&N}^weP{>U^CS}YtQ)vSlv0d(8iqVPZDa*DULa}HjYo%Yuiw^
zX6p9FHxu<@2g-IL<{SZ=sqVU?o^8C6p91T_b(!y6`WdASKju75sOd)>Kgu&;Hy(e3
zt(VtMg!m=x4CA|6cXT((F2di-kB~gh#TeN8p?;n(hWGY4uob;;rwDUm>OJ86JUkCq
zGuK&tmA(KEH}o&QGR9jw_9ETmgm(7Q%EQh+uy)kL&P!lRh8=Ul&dcE3&MR=WxTf0K
z4-Y%^FWNciQ7R76eUQ+OpI~{|nE=mHLj6K}cP7D>3_Iq8ohfi`XBw^+H&#0{@UTPw
zqMgH@vOTkOA11WpKLB~yc}-3U^{_Jswq)2b$8$H@sLk{j^RV2|>u|OF`l>l1Zkw(A
m94E|^XB!KAgFZiQ@il%EuIAqX9owho_#djNxq9}~;r{^uKp;2(

delta 2573
zcmZ9O+iz4=6vp?=v;_)9#{xsWR1mRZ@rtNeTdkGW8=#22p|w>hPl^{}VnR6j7Z9B%
z;~$_weK7N8eCL^GV~F8ZL{u!GQh&e6taN84S=rzB?RD1v&OU3OGvA&n|9Gn0Qz>s-
zl~O5nr`74`XFc<!^7rrl(OYcjOzFLa4;J1@DRrifbW0n*qmA!q<9plqzBc}98-G3H
z6RD$dZo?Oq(qiM|&Bwb+<;HLQ<K@0VY;<Agua1<KODXMyeOOLuu`#@LtfT3eJ4r*N
zO*85^=Q_^0j&rW#oa;E}I?lO{<)iG+N9EL5-hQ@J`m%9$U}Auw@4&Innaeo+lhE<i
zR8E%|ue8R?T-?p&f1_O_ekRQSomfhhbcNC1HEze8B|^>Xm+HS9>MepfYGMu9SHb6r
z&%wj3{kjHMKSPfD_YYX@I`|tjUHl7IKM8jA>tKJb?%aQbnyZVcT{Md=L(UMNw2Ct%
zwM{&rFi=AK2~CqenKHaqq3IaS8APjq)lalccYxJKS>-tMo#11P?-Q;ul8Q=GG%*(H
z0_SE{!PWXRGppfY#vNM&-+~$SsM8HrkHy!5V{vPautRHa1eh3Gzpno2WN*=fvW|$&
zUJq95!oq2cyV*Cu)nl`(V1I75b5_2IaI?iO_bdB1PmXP5zH{b;<z6ti-re-<CU2}?
znyD3=kT(*R@JP$|V8cGBfe$mP$DM8l$DLZz>(tDRrTW16F5e7Sv!ATv>bJt_!+K+?
zTHFSfTgyW!4-frdJ*bC=+ri<%+Fqw-Zg|)N&OK~}tGUhLVH;czW7E~*PO#isz6j*u
zVLNyZ59;CJE^v6Tw%4hd8y*J0xre*qYU?u(_rUdVbhcUyg5}oo<s%OdJHdKT4-fZ(
z!-KWGPR-o#a347La6eq_1Rl(HP7S=+?4b(?c@No3@;Ka%7Pn6yV9vhth42vApFOJn
zAfaX-ieuAv)rY5R)rV1bW$JF5apn<}-NX;AbMPqGO!co@Yd!{6cg@EcV@>rvgqmxL
zW6h!Z!fdVj1j<mR4%fe%sud$B!$hq4B-l)K?>p)_#w+<LupYcG^Ic0nPcw!eb4Cd@
z{fOhCjDg*F{Pj0qUe6%JlXNnSk2lZg1j;z!^YSAkkG^;o?D<f?z!$@F`yAMcp0|^P
zIkEKf;JhDRfUBA7s=i8JgohjES6&(Wn>%)Z>HUOu_A$!C&P!nJsE3`G!Ilg==7gPt
z;M~qDaJ4w5+BpOdJIt@xnY5LPDW)e0?f40nhn;Eg93|8*wsvO*Y{{@=PS}|R=XMUm
z)#Aiz=QViPVSdHVoLyGD!So!V9sdEy!_K^Olu!>l3t&rz9dqov=o`nd_a->E_ZD0&
sJIn`@q6OmcIm*w`a5kc26nvXGzjg6(J_1+s-+?#a+>u7-+*dpP2hy1%2><{9

diff --git a/piet-gpu/shader/gen/pathtag_reduce.dxil b/piet-gpu/shader/gen/pathtag_reduce.dxil
index d585c96acc50a24c7d0b7674f55413ef78e681b2..245c4922dfd27d0d2d5f30b89554bd048b453099 100644
GIT binary patch
delta 1033
zcmYL|UuaWT9LLZ7lbb)ix%t<$r`9e>o3vfjaTltwVj;ICv?;n}Hq_`gxM?@>&sJjR
zm=7JfH!(H1I4ZFsvWIlqEw~Qiq8R&h@03wS$GS1ba3LgRNajT9i<Gg4outinpU&s|
z`@;R4b2vr&O?#U+G<e{z8@;cu%}u9sf6kXQmD6er0ARn}_BBCW#C<`{LVAU`jz2!0
zK0!*lU$*7A_JZ6Vi0GA5tTO$3Qx6tM6SS@Wh{LtNb!mU$8>{=Cdx>4T<Gy$QW$efW
z`p1QpORv46d$Q|?OOul$EcN#C`di<Xo)oL!|MB<>@&k~+ulNM-rw6L8p;NwR4_Zgo
zqC10vH{?IKtOkGw3NVEOP-B41^hhp;^1%aA$wmWkGXJq^6jLSvJ&GpjC_a~ePcuwv
zA0aZ4Lpy7IQ(9*FqtnEQI|B&k{5e9izz)wzDl`IJlp4hN@}lhcOor22l*77ja&3Z~
z<Z!_X_r5=?J(;1x>^KjsDr#msd+Pt?=>WE^n8a#5Z$vU2f~juOqacFo)vfR>p!@;q
zhB}G`#4lTtBJJEb`%nTNJ)IgDj%PINs(~N?$FhI(+StBtJwr4NNLkv5?@>7oM8~Tv
z(33=0bIF?Z>2fi5pb5e&50oKsdscD68t=sKW&}<#CA6s6EYGas7-cSqSI0E$xKNY5
zyrU{ZGRj`pM}fI49+Q-=tX0<#I0h3L%|Jt7a=tmX??ppkIWekoRvH2mt%q&dL`yCf
ztP4yug|Ho2#pnMCEY<}w=DI*8b~R1~M%z6{BVhMkEc9IF6Ij&pb451>*SQW)IDuJy
zwKXFUJG?b`6mFi;CI{PuFcBN3_}>ILi@`nC^AiBY5)_Wc_zYgw^12yOBVi>TW`Vt|
ztpQ$#zD*zyY@V0b5~;M@dMWi4Z{`a&LFu~ou0=n?CuI9GE)-_f({KzP?6f4!RC@Ko
zIEr7$={F7SQAjxf7m1_4v7}ptz*mRJCioF-Uft}qbIOm9;X^A#&Ykr6ny$pxV6CgR
zf%S`6?o$WP7WGvEuDJ9ZTV$DC7?r0Jh7^QA6%)^xo|cW~Ts#H)5XyE_=j9I>%n$9H
z(SDmTFOpk>K2z5`!}+a_H;fLMacmG<xz~hhqQAF0Z0!Dt*<EPlK@+OUezxa~v8TfL
a{%Pdzm{1KH+1o1`d!I9Tl0vVl{{0IH%R~kM

delta 1148
zcmXw&e@Gi=7{}lHnM*uPOpIP6b;%l~rZGK(P0+RDPF>eVH&T|mAYHkn)%iy&TK7Xb
zNA51hsmn%PTr0YbTicavbX_zZUAhetbr76gBXuA{M5!w>=xCt~Ittsx#(!Sk&-eYl
z&+~nFd9PSjErw=aW0TvUed?VrJo(>~_rBAHKlvB}034q-e2d_h;ILQqXLJ#ngz-0{
zU1%!E#qrZwy!#<uVw@L@i>%)1{Ly)_ZpPudRn2<B<eguxw|3hf+&<%G>dbR}(*-9J
zd!tj@lD9JtkKfPEW)?Ol*D?z$SJA@;$APlsO!8(!a(X6Nojx`~-U-{)*KL{G4?eO-
zawWRmso&)q)qR(vzlP4W^j)5s(rP1Ynd894>6;C%TaKiB&zMMJAHwyDBK<odCv<FX
zDfZ2i6TOON-;q|J0)PV;;9&r^SsqLy&$gumps^{nW+nl6H}O948l=GJVIWg92!EP5
ztvZYrZpl6!dc>CVYoqC+`d~lec1ZwPa7i9j$g7M1up3tbkz0=ml@kKhpQmf-N(nLy
z<S7n0CAOST4+{`)qNua+D?8iXc#OrmdUoT4TtoA7`Vh_0lNzAE2DLE^c4p^b38MC}
zhDa3}WUw#VO}HbrezzDw087*Vm<mN}ig}qN0h&&2h4vGF5y-Au8fZ}Dp85|{teJdk
zJ_Q&J8S(Uoq|$>a)f!X7ah{^*IP9!cDhcE$u`I&4J_EIzQj%ED6rhxM9dBsSkY3<~
zZZy&!NW)lBL=Ox>@KsLj27@^@4rLW-2IFw10__-tyz44gfCnZ5KUy<5TaI4c9_6Xf
zPOBU-P?C>9p74aFQBO9Xe?ZWS4b5ymKTy0*SPfV{f0L0j$2!X5u+LLKQR8<De8IFW
zhJ9@cJr92`@hq0$O9?v0mu6uI*Z(A5m}6h57{_o=lw$O!KqvkPgf@owT`BGYONjY~
zkIki<{xEB~Z}#~$%5$1Fs`*mIkRBh4D+03o_m`B@G};M<3JKA``;mYT=T>;lI3&yB
zsUYZ74hbso;^4tOxD~kjLjeg+4Z3ph>NWsmqi|GE@uVys+I|#ZR`wDf{HwTU?c#ng
zXBCc#4$~+XhN>DHeSTg!Mzv8~z%s1ICrl#&S$5A#FlkNYmx$FEgGH+)CW=<uguOD?
zN$3grC732$lnmi*=y?VsGCX_j^PguF1JO5%O~K!)IC@oc9;Fr?R1>}G!CerZAQjr>
z7gj<Mstyoh`i5?At(t8qsaew2998SHXkMS!*;uvEVzh5)?S6ITXFK|`PPV06Ec+g5
w_w}mxuOoT?rcSm&tu`9c(SrIVDd#WGHF&I(eQax8ZLzL)Or1ap`7;9WKk?dYZvX%Q

diff --git a/piet-gpu/shader/gen/pathtag_reduce.hlsl b/piet-gpu/shader/gen/pathtag_reduce.hlsl
index 291243e..1332429 100644
--- a/piet-gpu/shader/gen/pathtag_reduce.hlsl
+++ b/piet-gpu/shader/gen/pathtag_reduce.hlsl
@@ -38,8 +38,8 @@ static const uint3 gl_WorkGroupSize = uint3(128u, 1u, 1u);
 
 ByteAddressBuffer _139 : register(t1, space0);
 ByteAddressBuffer _150 : register(t2, space0);
-RWByteAddressBuffer _238 : register(u3, space0);
-RWByteAddressBuffer _258 : register(u0, space0);
+RWByteAddressBuffer _237 : register(u3, space0);
+RWByteAddressBuffer _257 : register(u0, space0);
 
 static uint3 gl_WorkGroupID;
 static uint3 gl_LocalInvocationID;
@@ -82,12 +82,12 @@ TagMonoid combine_tag_monoid(TagMonoid a, TagMonoid b)
 
 void comp_main()
 {
-    uint ix = gl_GlobalInvocationID.x * 4u;
+    uint ix = gl_GlobalInvocationID.x * 2u;
     uint scene_ix = (_139.Load(64) >> uint(2)) + ix;
     uint tag_word = _150.Load(scene_ix * 4 + 0);
     uint param = tag_word;
     TagMonoid agg = reduce_tag(param);
-    for (uint i = 1u; i < 4u; i++)
+    for (uint i = 1u; i < 2u; i++)
     {
         tag_word = _150.Load((scene_ix + i) * 4 + 0);
         uint param_1 = tag_word;
@@ -111,11 +111,11 @@ void comp_main()
     }
     if (gl_LocalInvocationID.x == 0u)
     {
-        _238.Store(gl_WorkGroupID.x * 20 + 0, agg.trans_ix);
-        _238.Store(gl_WorkGroupID.x * 20 + 4, agg.linewidth_ix);
-        _238.Store(gl_WorkGroupID.x * 20 + 8, agg.pathseg_ix);
-        _238.Store(gl_WorkGroupID.x * 20 + 12, agg.path_ix);
-        _238.Store(gl_WorkGroupID.x * 20 + 16, agg.pathseg_offset);
+        _237.Store(gl_WorkGroupID.x * 20 + 0, agg.trans_ix);
+        _237.Store(gl_WorkGroupID.x * 20 + 4, agg.linewidth_ix);
+        _237.Store(gl_WorkGroupID.x * 20 + 8, agg.pathseg_ix);
+        _237.Store(gl_WorkGroupID.x * 20 + 12, agg.path_ix);
+        _237.Store(gl_WorkGroupID.x * 20 + 16, agg.pathseg_offset);
     }
 }
 
diff --git a/piet-gpu/shader/gen/pathtag_reduce.msl b/piet-gpu/shader/gen/pathtag_reduce.msl
index e82577c..6c0a64f 100644
--- a/piet-gpu/shader/gen/pathtag_reduce.msl
+++ b/piet-gpu/shader/gen/pathtag_reduce.msl
@@ -103,15 +103,15 @@ TagMonoid combine_tag_monoid(thread const TagMonoid& a, thread const TagMonoid&
     return c;
 }
 
-kernel void main0(const device ConfigBuf& _139 [[buffer(1)]], const device SceneBuf& _150 [[buffer(2)]], device OutBuf& _238 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
+kernel void main0(const device ConfigBuf& _139 [[buffer(1)]], const device SceneBuf& _150 [[buffer(2)]], device OutBuf& _237 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
 {
     threadgroup TagMonoid sh_scratch[128];
-    uint ix = gl_GlobalInvocationID.x * 4u;
+    uint ix = gl_GlobalInvocationID.x * 2u;
     uint scene_ix = (_139.conf.pathtag_offset >> uint(2)) + ix;
     uint tag_word = _150.scene[scene_ix];
     uint param = tag_word;
     TagMonoid agg = reduce_tag(param);
-    for (uint i = 1u; i < 4u; i++)
+    for (uint i = 1u; i < 2u; i++)
     {
         tag_word = _150.scene[scene_ix + i];
         uint param_1 = tag_word;
@@ -135,11 +135,11 @@ kernel void main0(const device ConfigBuf& _139 [[buffer(1)]], const device Scene
     }
     if (gl_LocalInvocationID.x == 0u)
     {
-        _238.outbuf[gl_WorkGroupID.x].trans_ix = agg.trans_ix;
-        _238.outbuf[gl_WorkGroupID.x].linewidth_ix = agg.linewidth_ix;
-        _238.outbuf[gl_WorkGroupID.x].pathseg_ix = agg.pathseg_ix;
-        _238.outbuf[gl_WorkGroupID.x].path_ix = agg.path_ix;
-        _238.outbuf[gl_WorkGroupID.x].pathseg_offset = agg.pathseg_offset;
+        _237.outbuf[gl_WorkGroupID.x].trans_ix = agg.trans_ix;
+        _237.outbuf[gl_WorkGroupID.x].linewidth_ix = agg.linewidth_ix;
+        _237.outbuf[gl_WorkGroupID.x].pathseg_ix = agg.pathseg_ix;
+        _237.outbuf[gl_WorkGroupID.x].path_ix = agg.path_ix;
+        _237.outbuf[gl_WorkGroupID.x].pathseg_offset = agg.pathseg_offset;
     }
 }
 
diff --git a/piet-gpu/shader/gen/pathtag_reduce.spv b/piet-gpu/shader/gen/pathtag_reduce.spv
index 6dc35b845e6222173a4d994a2e1bc950f1404231..9fc105fcca998a019fe17a6483ef5775a4eab006 100644
GIT binary patch
delta 1744
zcmZ9MOGs2v7{||@ab|SV3Z}ux8d7l+ZG<)v3PxZpL|VBB5reUyWMmo_F6xfZ(v8e6
zf>s3uK`{{IA`DvevX{N>q1i*TOndpt?Ejm&r_2NY`~AP?`Mz`Ry(epWYD)aR%8)UR
z37bOW?=5L`CNDPc3sm1X#w1b)<4MVYDKO6@J6e;i?VdjLpB^ydy(1B4JhK=m)w9GY
z)3Z7VuXC_B2g}?Er8IV?c6PQI6I4-jR-!YtA6vlq&F8#HJZL7x4;*f3N+b`}Clj3=
zdv@hpS(3dW-kwNyQ7357S?%%mmPA|Ik$CE#-SO=4WHOdu7B(iqmv`ifu(6YkhMIhI
z9irdvqDXYy3#*hY6m~0Bo=~I;Pl14EtxzxRqhwH6ClEQ287Yc1Fr!}}g4rqrmetF}
zt)F>U2^&2BZ%Hb$6BU<oy)DsOogKmUR<A19Y%NEFFW(EDy;$t+%bpD{IF2uSy12<H
zDwH>7LYi(Q9LY?D3*1er#6(OQ6K4+e&C2~G$?;*^;pxnV4gHuTv#4k;{;MQDYEjqD
z74xow83A_cE3qP%z(5qF5(7|Vt3r=T`Z5En2g`X&_4ekXxFX)t<@!+e>YA+IJuWl2
zF4Sq)Zpj-$pO6-6v<V&D6tG+UmUzE#O*!~)3#SC^Y0owwP9DAsvYnB<BVfNI%ox9G
z?uoGGwGqy{0=8Mr1h-8+5MSBUplsM}QxC=a1>4jkahk$zn|drx9=<x`2Fw+aC&Fa`
zC&f6Zrvmm`-3y0_Zhg<iPk25K7Tx+@h!am9zXJ5_IwNYn6p96M@YT8sq7aTbO>8Ge
z4yzHa1}wjxSJL~0v|uNHEsotbF)U79PLBac1O_CI#?Xg_Oh<Hm)f<t1!5-15IJS)f
z!Tpk>ne)*y_pR8<t>4Lp{hW8h@5QmD1>LY2lVowh_&x~OxL~-w;4k8Q8)?CgH6@PS
z8_PA*Qu1>J4Ij@pfe$Py5WLrW^}dS}%VL8b$sYpo#4;8-aW?0tI63;9{^rbxlaoIR
z*RP3;A<!nxij%|#Xg8V@Cy7508=(_tbLPd#;g4i<7R1TfB9MbloXz<qKD|wk0>LIN
Wij!0=kc3VkZ<0OYv20gqW#|taI2En{

delta 1789
zcmZvcOGs2v7{||@r<ztU4OW(9gPUq2L=d511lA&?m5UHF7%NIfCSA0sJMP`c?4oE@
zP^)AhNGM#i=>4*HR$BJZEX^LivikpK?n#p_{OA0?|MxiGch0@jQF)@W$nT4Vjd4tY
z$uoh@qFQIPtA2N|;=VB^(YimL)EYDa^Gs_?W3sW?v-|$BgJ!sMAmR+CXM@H1Epkfq
zTM>d+JLvVH5_d!?wQa3i+nS6CsVFv2qOEl=TG05-=bR)SGNa=A4mH#zk_T&&iMEy<
z+j6a_S+g_VoJby_PKf5SPjkGvA<@)yINmy^<3f)olcfn}VPg_}IY+h#8~ZZRaGj5?
z!}QypABhetgs_;9Cu~=$JmE+=o`C0Bs<lSeN3H6`EE9+vNe|>lvNQX=2xeO=@Q{IB
zCRji7E)rG>_<w7qB0Evh&PhK!R%1rA`Eg#~o3f49u!W-RMn3WRjbN8iQ4bsYGiM8C
z94C-DU0CO=pwxP)v~o?zQ*@~b>nH{aBI&V$fV*0i8J}Y*73X5H*C^|gRw|%v@oWZ1
z!#<{!!BsRH|5YnKYEjo7LaBEEbOF$*?>HKT1S(Kaf?*)E4GKM~)t|n!xVMZ~TCaLG
ziu>f1o^SVMx|U@8?p4WvS*x{cyVmPMw~!JlwFw>E5YVlCQ@lsGq8$9Ugi``~%F}L(
zlZWrTv@=@o2<R7t3FCLoJrSO1ZG`i#fHtX_;I^q=@r6x2kcMuXdMMr_*rpzd(-gXG
z>ajR^_^OQ?G?zu52p0vM6yu<t3g}h37Y-BK`ksp)_k0{Iw)MRbC!Re1H?TMBjHvli
zC=|%SSLG^*LO7;0v7MZ?Uyc0ldS1!bEu;iHdA~TiZDK&2x||*Zz7`m;U!XB;V$yq}
z%gf)0^a%Ed2F1}<2?Y0O9Za8$mbh=l7H<7c8u~f!hTn^$r3Bru8Pdwb1>^f5pmD))
zd%<Jkd>bjjj`c+x-5bj_<1%t{g$y6hH-Qf<DiFNOd-cAH6U)N}9a?_~#1qR{*u>eK
z32}1tIsMI<6elNl6s})gGu2$ulsHLzfOey4agt&J8(|Y?H<}SAhkq!W^HZFhbpkoq
j#Mzu#@$rrN2?#dnmpDlk0!i2e@+R3M9?Bdjj)ngK3E2%Y

diff --git a/piet-gpu/shader/gen/pathtag_root.dxil b/piet-gpu/shader/gen/pathtag_root.dxil
index 1f27f265c5be01104cb77bab185d0e73358e5fd2..77f12e6db0ec75e4d05be227c033a9793a2cc1a0 100644
GIT binary patch
delta 251
zcmV<X00jT+B<v&<L|8&Y_OtYGzILa)<opDnZLM_Okrasm0kNJf0RaND%mIx79Y|2X
z0)<k5P-RdGr7uuSMM1bgsEq*7odKXQ0|2vj13Cl{24C8Z1=6d-9}(NwC5Xy;lfVb}
z0RofE0d2D-2`B*pNV84~sRIWJqd*EYpd^$24Sy7LJ7RzUAP5E#9EgDwjKOF^02K3+
zkq)^Q5<Gwr0|Wp;Fo;&6585E`gVBTlC`Pkj4-5eWgD><#lW-7l15pq3lj9I36W9kZ
zm}M}AV8p-xKoAU&mbeFdlN=G~8G}#2h(QB@AQ((=0Twe6ki|h3WI+}*QG}Bn5>E_C
BS~~y$

delta 235
zcmV<H02Kf1B<v&<L|8&Y8iY>|o5&7LaVf&--J{<Dkrasm0<oSg0RaTF%mIx735Za@
z0)<k5paGL%0~WJ*12_Z_s`D(i4-As4vo}(se@E)nlfVb}0R)rG0d2D-2`B*ph_g-#
zsRIW}qA&`iKm(Kh4Sy7HJ7RzUAP5E#9EgDwjKOF^02Jeskq)^Qk~@GA0|Wp;Fo;&6
z585E`gVBTlC?>OD4-5eWLoa|slW-7l13?e>lj9I36Tk;Bm}M}AV8p-xKoAU&mavC>
llN=G~8ADINh(QB@AQ((=0Twe6ki|h3WI+}*QACp+5>E_0QIY@v

diff --git a/piet-gpu/shader/gen/pathtag_root.hlsl b/piet-gpu/shader/gen/pathtag_root.hlsl
index f1ec389..7ad806c 100644
--- a/piet-gpu/shader/gen/pathtag_root.hlsl
+++ b/piet-gpu/shader/gen/pathtag_root.hlsl
@@ -7,7 +7,7 @@ struct TagMonoid
     uint pathseg_offset;
 };
 
-static const uint3 gl_WorkGroupSize = uint3(512u, 1u, 1u);
+static const uint3 gl_WorkGroupSize = uint3(256u, 1u, 1u);
 
 static const TagMonoid _18 = { 0u, 0u, 0u, 0u, 0u };
 
@@ -21,7 +21,7 @@ struct SPIRV_Cross_Input
     uint3 gl_GlobalInvocationID : SV_DispatchThreadID;
 };
 
-groupshared TagMonoid sh_scratch[512];
+groupshared TagMonoid sh_scratch[256];
 
 TagMonoid combine_tag_monoid(TagMonoid a, TagMonoid b)
 {
@@ -73,7 +73,7 @@ void comp_main()
     }
     TagMonoid agg = local[7];
     sh_scratch[gl_LocalInvocationID.x] = agg;
-    for (uint i_1 = 0u; i_1 < 9u; i_1++)
+    for (uint i_1 = 0u; i_1 < 8u; i_1++)
     {
         GroupMemoryBarrierWithGroupSync();
         if (gl_LocalInvocationID.x >= (1u << i_1))
@@ -97,16 +97,16 @@ void comp_main()
         TagMonoid param_4 = row;
         TagMonoid param_5 = local[i_2];
         TagMonoid m = combine_tag_monoid(param_4, param_5);
-        uint _211 = ix + i_2;
-        _78.Store(_211 * 20 + 0, m.trans_ix);
-        _78.Store(_211 * 20 + 4, m.linewidth_ix);
-        _78.Store(_211 * 20 + 8, m.pathseg_ix);
-        _78.Store(_211 * 20 + 12, m.path_ix);
-        _78.Store(_211 * 20 + 16, m.pathseg_offset);
+        uint _210 = ix + i_2;
+        _78.Store(_210 * 20 + 0, m.trans_ix);
+        _78.Store(_210 * 20 + 4, m.linewidth_ix);
+        _78.Store(_210 * 20 + 8, m.pathseg_ix);
+        _78.Store(_210 * 20 + 12, m.path_ix);
+        _78.Store(_210 * 20 + 16, m.pathseg_offset);
     }
 }
 
-[numthreads(512, 1, 1)]
+[numthreads(256, 1, 1)]
 void main(SPIRV_Cross_Input stage_input)
 {
     gl_LocalInvocationID = stage_input.gl_LocalInvocationID;
diff --git a/piet-gpu/shader/gen/pathtag_root.msl b/piet-gpu/shader/gen/pathtag_root.msl
index 923e77c..65e3741 100644
--- a/piet-gpu/shader/gen/pathtag_root.msl
+++ b/piet-gpu/shader/gen/pathtag_root.msl
@@ -67,7 +67,7 @@ struct DataBuf
     TagMonoid_1 data[1];
 };
 
-constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(512u, 1u, 1u);
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(256u, 1u, 1u);
 
 static inline __attribute__((always_inline))
 TagMonoid combine_tag_monoid(thread const TagMonoid& a, thread const TagMonoid& b)
@@ -89,7 +89,7 @@ TagMonoid tag_monoid_identity()
 
 kernel void main0(device DataBuf& _78 [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
 {
-    threadgroup TagMonoid sh_scratch[512];
+    threadgroup TagMonoid sh_scratch[256];
     uint ix = gl_GlobalInvocationID.x * 8u;
     spvUnsafeArray<TagMonoid, 8> local;
     local[0].trans_ix = _78.data[ix].trans_ix;
@@ -111,7 +111,7 @@ kernel void main0(device DataBuf& _78 [[buffer(0)]], uint3 gl_GlobalInvocationID
     }
     TagMonoid agg = local[7];
     sh_scratch[gl_LocalInvocationID.x] = agg;
-    for (uint i_1 = 0u; i_1 < 9u; i_1++)
+    for (uint i_1 = 0u; i_1 < 8u; i_1++)
     {
         threadgroup_barrier(mem_flags::mem_threadgroup);
         if (gl_LocalInvocationID.x >= (1u << i_1))
@@ -135,12 +135,12 @@ kernel void main0(device DataBuf& _78 [[buffer(0)]], uint3 gl_GlobalInvocationID
         TagMonoid param_4 = row;
         TagMonoid param_5 = local[i_2];
         TagMonoid m = combine_tag_monoid(param_4, param_5);
-        uint _211 = ix + i_2;
-        _78.data[_211].trans_ix = m.trans_ix;
-        _78.data[_211].linewidth_ix = m.linewidth_ix;
-        _78.data[_211].pathseg_ix = m.pathseg_ix;
-        _78.data[_211].path_ix = m.path_ix;
-        _78.data[_211].pathseg_offset = m.pathseg_offset;
+        uint _210 = ix + i_2;
+        _78.data[_210].trans_ix = m.trans_ix;
+        _78.data[_210].linewidth_ix = m.linewidth_ix;
+        _78.data[_210].pathseg_ix = m.pathseg_ix;
+        _78.data[_210].path_ix = m.path_ix;
+        _78.data[_210].pathseg_offset = m.pathseg_offset;
     }
 }
 
diff --git a/piet-gpu/shader/gen/pathtag_root.spv b/piet-gpu/shader/gen/pathtag_root.spv
index 88e20b907c6f45e137a705f3462f2a62d9151c8c..3783b49cb1351c02c3e5aa12cbfd4f26c93e4936 100644
GIT binary patch
delta 1869
zcmZ{kIcQW-6o&82B$=oLl@y8=3W^|_La-50iIBoZv=u>8L{tzZ2$m!16e8+~`x5uv
zB#!$!F1RM{`x>>fwQy_n`^>vK4*1~b{Ac;^J?GxbdtKL6SCdJP$cB&#eM5ElxVmO`
zD%H7QNTz8=2w~BZc}>mO%1{vwU>7$wH!e*4L;v`ha2&sR(NbzNp$MN3<TtTVa~oc|
zMpyn*=40wZp}S&aI+e<Qu9`J&10`uNhApeG8$op{gyHy4lbrFfWT#+1<Xf}5ayH8l
zVIo!iur^!j+ia<Cv(X<6FP8Y~Z93-lZM@X)O8k08oEk!*y?S(IZau+upaJ}#5&bnF
z556&!yob0W3P(#82%F$5z(6pG8b{m=j7NV9d=<FHrslVTIl%bQP?-yJA=`jLlV3Uh
zHmsiwncEK9fDOrG13Tg64eTOjJT|Z!UfRGOd>b&{5^}v=?Zr`O(#2O|_krcWoQrvs
z`+@PvH0kx}$J{}<F>_v}e#{+)w<I~QReu8Kk<CGlfc`*{Nf)r#_;)65!#cV4e^+so
zT=_|~6Ehy0=zzO!9Iz9fb3zN<)PO~9x)$rES{Lx{jscHyS#p<!BsL7I-kdS@tFh{R
zN99dVNKW$v`-eUhDE4&|_;a@>;pG#ZB4+$B4ft+O17kiw-_2R9pYO)p8DPvt<ni5{
zgZHk@`N!$b6ExA1Om_ip%*PUs>>}I~GZ&Bb65MzkLf;dJ8*9xKu{LV%GS;Rp4}g53
z+pE}XAmm$Xhvu$Bil8Hz^9I~l18_F|D?pXY$m1Bd;KjtxV($RY))sCjzT^0Z=myh(
znmEo~FU5oce;EXN-zxSV@Tt0xIcxmw7>^Cy{~iD9Kfo`7_-Y=)tuY=qeFS%%dfwR#
z_79J7rh`F1!OY}Rp1>{ka|-{i&w%yT+K|3=QS%(GCIzDA1zgQAphn-i-x~H0FM*;o
zh@w|;MWcWsee2>zui<JkAZp&g)zkqs`qo9wTezAE5H(Hja1@OPiu5h?>G(>#b8S5G
b9(WPNBYqF>0dcw_+*s+;yIOct`!oFuuloSY

delta 1885
zcmZ{lNo>tg6vpp+uM;9oV?kJu5Q)$oi$)BUSl9?#k%$FBB9%sBB~Q9zA?k^ln#Zc)
zndee-Ypi)z8(Rz0hTre?fAwkNC11~XzH`pK=iGbWYkT#Z>i(JZ@bVB+p>L=NpBwwn
zPNk188l0K8D}+$DbiurOY-uP7hp<cL*3Vs>_|5<LnQ#ifzHS-0na~PfSjfwSJJ?ur
z4_>@RNA63eE7g_nC>fDXrE_1)X4P&Wr4Ja*mSxyYAYJ5dhF7FQ7)Jbg5;JcM*^{tc
zx#i`%vo_CAVLVy=ur^=p+kCNa^U)s!cYx?y?-0elgB1H6iC@e3Q$ol$Rg5alZbV!U
z9O)O0=&uDi@Pnb`AJ839I9jqm*aBY#s=x$t9C0g9kN!6J8qmh3#<znxK)ofDX2X0)
z6R^<Wcdlk9){hMt+X2*UNFE#54e#B+9yImXz+QN91N-o8K)oSkySv(tW1&G8UyVHg
zRsv%#=20F5>Jw?w>(h_1!*DfYUZsAF9f3C_F|Sp>hVv-TLRvs2u*jebn5_PjiR0wj
z|6RpVV!cn|7@B%);yB!O<A5jNStm5nO%+UX({)&nTk8VeT^@LpE0ViBnP@|?)*Dl^
zz7cD^@2GdvHOXn7X8*$X$YNhN!r!|+1MfZ2Sv2)#8t~nm18P1%-^~TAAK#6!^FYl;
z<ni5Hgm<sa`N!!lAsT2%rn?MR^RdJuy8`#bjKyQU3RjOq=z9WjW391Ptc@DGhPA0{
z>>s|-Z5#Fm@MFSFtO?rK;4Qd|980FZ4Og2AoLK)la4l=)vGy+9+SqhE+!N2*nLiW?
zPT3E8A4~%l#exTkZoWSx`h^#aeF!Rn>licV4pu!j@Tj2o_>b{vDEs#@JwY(ZK-{zw
z?n1RZzZux4U^*BG%%7QD)ib!cejMUC_60ECTpQ9iFV?(-cd|eKRC<bD!7UmJEYdeI
z7QKeMQ5wXJ-oULH39QjKFV?(;Tay8?<{jLc>dE}ENWsKd^d4?e35Z1>;1-Pq*66>_
f&8!+S?jy|GjbF@9aJ4vwF}32S*Ou=X@H_nnD5wHn

diff --git a/piet-gpu/shader/gen/tile_alloc.hlsl b/piet-gpu/shader/gen/tile_alloc.hlsl
index 010e714..5231c1d 100644
--- a/piet-gpu/shader/gen/tile_alloc.hlsl
+++ b/piet-gpu/shader/gen/tile_alloc.hlsl
@@ -261,7 +261,7 @@ void comp_main()
     for (uint i = 0u; i < 8u; i++)
     {
         GroupMemoryBarrierWithGroupSync();
-        if (th_ix >= uint(1 << int(i)))
+        if (th_ix >= (1u << i))
         {
             total_tile_count += sh_tile_count[th_ix - (1u << i)];
         }
@@ -271,46 +271,46 @@ void comp_main()
     if (th_ix == 255u)
     {
         uint param_4 = total_tile_count * 8u;
-        MallocResult _477 = malloc(param_4);
-        sh_tile_alloc = _477;
+        MallocResult _476 = malloc(param_4);
+        sh_tile_alloc = _476;
     }
     GroupMemoryBarrierWithGroupSync();
     MallocResult alloc_start = sh_tile_alloc;
-    bool _488;
+    bool _487;
     if (!alloc_start.failed)
     {
-        _488 = _92.Load(4) != 0u;
+        _487 = _92.Load(4) != 0u;
     }
     else
     {
-        _488 = alloc_start.failed;
+        _487 = alloc_start.failed;
     }
-    if (_488)
+    if (_487)
     {
         return;
     }
     if (element_ix < _305.Load(0))
     {
-        uint _501;
+        uint _500;
         if (th_ix > 0u)
         {
-            _501 = sh_tile_count[th_ix - 1u];
+            _500 = sh_tile_count[th_ix - 1u];
         }
         else
         {
-            _501 = 0u;
+            _500 = 0u;
         }
-        uint tile_subix = _501;
+        uint tile_subix = _500;
         Alloc param_5 = alloc_start.alloc;
         uint param_6 = 8u * tile_subix;
         uint param_7 = 8u * tile_count;
         Alloc tiles_alloc = slice_mem(param_5, param_6, param_7);
-        TileRef _523 = { tiles_alloc.offset };
-        path.tiles = _523;
-        Alloc _528;
-        _528.offset = _305.Load(16);
+        TileRef _522 = { tiles_alloc.offset };
+        path.tiles = _522;
+        Alloc _527;
+        _527.offset = _305.Load(16);
         Alloc param_8;
-        param_8.offset = _528.offset;
+        param_8.offset = _527.offset;
         PathRef param_9 = path_ref;
         Path param_10 = path;
         Path_write(param_8, param_9, param_10);
diff --git a/piet-gpu/shader/gen/tile_alloc.msl b/piet-gpu/shader/gen/tile_alloc.msl
index 3906536..49bd1c4 100644
--- a/piet-gpu/shader/gen/tile_alloc.msl
+++ b/piet-gpu/shader/gen/tile_alloc.msl
@@ -272,7 +272,7 @@ kernel void main0(constant uint* spvBufferSizeConstants [[buffer(25)]], device M
     for (uint i = 0u; i < 8u; i++)
     {
         threadgroup_barrier(mem_flags::mem_threadgroup);
-        if (th_ix >= uint(1 << int(i)))
+        if (th_ix >= (1u << i))
         {
             total_tile_count += sh_tile_count[th_ix - (1u << i)];
         }
@@ -282,36 +282,36 @@ kernel void main0(constant uint* spvBufferSizeConstants [[buffer(25)]], device M
     if (th_ix == 255u)
     {
         uint param_4 = total_tile_count * 8u;
-        MallocResult _477 = malloc(param_4, v_92, v_92BufferSize);
-        sh_tile_alloc = _477;
+        MallocResult _476 = malloc(param_4, v_92, v_92BufferSize);
+        sh_tile_alloc = _476;
     }
     threadgroup_barrier(mem_flags::mem_threadgroup);
     MallocResult alloc_start = sh_tile_alloc;
-    bool _488;
+    bool _487;
     if (!alloc_start.failed)
     {
-        _488 = v_92.mem_error != 0u;
+        _487 = v_92.mem_error != 0u;
     }
     else
     {
-        _488 = alloc_start.failed;
+        _487 = alloc_start.failed;
     }
-    if (_488)
+    if (_487)
     {
         return;
     }
     if (element_ix < _305.conf.n_elements)
     {
-        uint _501;
+        uint _500;
         if (th_ix > 0u)
         {
-            _501 = sh_tile_count[th_ix - 1u];
+            _500 = sh_tile_count[th_ix - 1u];
         }
         else
         {
-            _501 = 0u;
+            _500 = 0u;
         }
-        uint tile_subix = _501;
+        uint tile_subix = _500;
         Alloc param_5 = alloc_start.alloc;
         uint param_6 = 8u * tile_subix;
         uint param_7 = 8u * tile_count;
diff --git a/piet-gpu/shader/gen/tile_alloc.spv b/piet-gpu/shader/gen/tile_alloc.spv
index d4a6e31710c5991852f7811fe336203f50b6c276..55d62adefde50e2cd91a9a6a32d822c8d0f4267e 100644
GIT binary patch
delta 2680
zcmZ9NTWnQD6o%)V_LR1^Kq*jb6<XSAr3FQ;h8RqmXevG!i;qg$6GGC)Y7MQ#7qivq
zgOB7q`>tY$O=Dsri4{=sR;?FA5yiVkyr5DpTJiffd+^Ym%vt~cuQhAdtUde8!JRMd
ztgpzoFA5<SYD0CnakBoVx=n|3A&d_XPK=ChA1scJj%^De6hdV<g8pD}NAdoY7s63^
z>K`1RDDIdbrVvh)9;<8YoX6wz@oc`c^ZD}Gm7!32wr+J&zd&8KfG@v@7eW(!Vr-%~
zI=F4@p$E>dc{y5I!QjZlltvX5p%rZ;t9#u=IV!`o|MOCReZ$mn{ct|lJ2kuHp2EO~
zRDO<mtyKI8xDwpJxH;qYMos@E<e_z;pV18;z+Q^x{GL!1))3n}^~j35b46QcsA+7c
zt&3cHVgdiFjIYFRLu0($W=5_pl-inF5<NCFdTxW;(8xE#ZD{0L8(Pk7W;wT!k(=LU
zMy|D)885xuR6ljNskJt#!tp944`HXVdw3_+;C<j%82V1}mJVQj&e9fr`x$DdLq5fu
z(9QK4S{ozs*TF9`4v<&>8{i?V`s1OHq~bw{6Lt0xIEj6ev4^2ACVMN>BYzt{#dtJJ
z{tj5pxLYV%pV}FUjKyk?QS3BDOo!4;+v47Lk;-eHfjc#p@*dbyj5Do#rnI@erSpA|
zyNd<Pg4MQ@+(>=1rKj3kk`F+x@m$v5kHBi4V*7gt>tlcQeauj^zv8&%Pr=sqM7HJ6
zz-oILuITYq+lb$<;dAiv5LTrQ1N<x$?MYnVP^qt@CHVqXuTh?Qxh<$6&6i*`H)bgw
zkz;W6XEOd3Sk3KxgY|Jc>R&U|+>Y41Tk(Gjz86d0(vgmu<U6?9Zv3wCJy@*>w&IvW
zdy>;^%nxAKbvMSF-x_|x`dEW8KQh#;K|CLFDOStcyDa<+>0<P8FfH(SX=`&+@(apw
z#tT_lzk-cacd6g8J}#y2H-?%^iR0V*1MDeyKFjkbSgpuh^UQ&L%%krl<L_KZ(nOM&
z=oHvQd_w6LbQ-KS#N>Ev&%iyl`r^E^V0GWI|6D1jNEfzP7Q5gplFOdY<&yJGgJPn8
zz$P$}e~I$w`xmSh&*(YuUWR&%A#ZrPEG7pxM%@kg2UY?0&~0J;bYd;G5>bZ>``;=T
z7qBiQt^%8IBPYZ;i@@Dj<BijA+&uC6T+n~UYOvZCa4fV2Ezxlo@1!3`Em+Y>-Ya#C
zT}+DKyxbU1&>iTD!5gxq_hpIK!Iz-x8^~VaQm~p&ET{p_HIbwhS+RnFR(C6bjo==V
z>GS@O>$@G_3b2Lg^DdFw5lmRg&z)HBAm{1#UB^SKwTMY<bDnK`ki0QtTP_1PGvYgL
z0jo`u)CIk))a-(-y#jo+RBUccR)J*uvVE@xtL@5WdabF&mD<1_b+fjL8J4;R>yg#Y
zvE>eIC&Neom6=}~b6sUGO%jDJu%*Ti`)V*jLEVo~-p(+&i;3ghwP1^lyc=xpcp3D7
zjZwFk9_jUY&T@LDAx`voUPHiohCB4qmW#cxO(NcaEfO1F;C1j;xbep6ch`2nLbb85
c>%kWa>tfOk2#UMQ+NxIG){dzs*Y02XAK5Ka`2YX_

delta 2668
zcmZ9N+iz4=6vp>VJEg6yP)ey~6k6I!rA2HNNMkS|(V#vUiw{0%QznGO#%c{ML=(+X
z<3B*=*>@v`*i1|`F_D`JR>k54uXw*kyrAW#h`-<TY#e)%J?s0vwf5R;ud~nW7k3@r
z)mWA7To6Jg)Q8%zVXARM!^V#?AxzviGC4Z7eWWlpHoh%{kPFq}Q}l-lI|>h!c`lrQ
zm;ECXlZ73V#N@)M;u8(cE&t$g`b0Y4*;U1t8df9=vXwci!{zgM^E}@2Kb{L$z$eEi
z3u7bO#&<qAx8_x7<qAedAEq>_s0tluqiNmO{+FY8cVkn|Y6eQfjT70-`qIq8dvZhX
zQSU{f+wpt|t_C+TZcMoyQPbZ*HX0T-GkW1e*fuoh_l26UlGyd7$6D^n4B0S4O=A;n
zSmfGM^Z0B9Uy9#0#&~&EvA;E+sMy9zk8O<n7W}p`^4s9HF><YKtmL+^lH0z>UD+0<
zyg1$3Sm%!91ErbPRrQHI*iW(BvHP%(;;sQ70IToioz%h)f?r{%pW!XNiuE~9yuL#W
zwX-2x#v9W-uc37?B7YtH0^<;QStkjHAzMkL<H<QO905B~-zKyv>>G@|41F=#bgD=G
zCcMOWJk9eKSk1VbC|aM|If|SswhB<(-;rPUHgaX1N8yIY<=+9jym6+IA1w}c=DTM=
z?kFyL46J5HjnsFn_(*3yc^Bjw&r+T}eIKmm8Mdd#u|D=x-v<medn%6G{1|L)Po>-Z
z1g!Ql!xcTeYJ>Rw7Cr+n31NBJVSpcnqJ4-99542F<&)1*^%~`wm)n9m(tH6{b7Pj`
z7QTe5Ka=vWz-n&i8?2ApQU98u=61y99mfAH_+Bh|i_=|o$#-zIJ@{SYd$3voY{fB$
zc9yej%nxAKbvMSF-x_|x`dEW8KQh#;L3|-(%2+LJ@8a+?q=&JI6Ka7ci^FZL$uB4;
z8PBI>{R%cx-KBoV`nZ(7-xz8xC5~_J53r};xirt8V6_5s&2t9qV;+5{8GmQ^QcNU?
ziOzye6u+8tV70AGj>q;q++(XR&if0j?mM2v%I&KQy9SG0@)gK6=S$|C2E{}dz$P$}
ze~9wvy9ic`XOz6<^Qj(VGH6d%#AMSLbvNMOS2fr}w}th~6YH@#L>(^77hLAz8n6qA
z{hu}AASc8*wcuW?@y6*l&VOosF6jSa9awD(I2Kxumgu;Xcd{ABWne`od8;%sb~7n{
z^KxT6LARqX1ox*&?@tr2fiFVWH$)!)b(_FyKCz(1a4w0YT#*$k7-)4j6SxH2M>2ih
z9&&xR;#&%~Fn!)6ayx<vt^C}9^#*dDe&2OGv|5Xp!Y<3QZ4Z+7Wo%0xyo?dw@s(h;
zeI#{3uPQaWU~89yPZSGn&B+Ro>_EEjcCgy+bf%Y@T3l%**rRS%FF3iSc49rU+S6>g
z8+$dwNB^qSuZ_8S?4?Pfa1Gc}<A=Q(Oi)nwBb0YBOzvXhIJXyUv61(I%^fd;HDF`Z
z?WIS0ZI-iKKGP5<dOWWqU@gNPdS%PS{aBNT2e1WV;|shV-VQh3IQ{P04p^u*7Ip*p
ZQeizzT8E&xr=qQDm2K@V?Oc6u(Z5%fP}cwe

diff --git a/piet-gpu/shader/gen/transform_leaf.dxil b/piet-gpu/shader/gen/transform_leaf.dxil
index 102d2f06e3d65390fc3cbca7f118a929ac59002b..32ec39935ac137a362c71d89ded97dd979d0760c 100644
GIT binary patch
delta 531
zcmZ3Wvp`40CBn(M$*{lV*ag0LZ>gHSVmlYUnkdrA$hdJ{F(XF@1DjxzA|vN!Nv55Q
z96}-w943k=L`=TTEIIizb3LQ<=3<t!%mPBD+gx9&SFO+s^I3n}EwOlV0LM2DC5C2>
zgA*E9Hv4fNVPu-nuvvsVnORU^sYF20)Bp+AAPv^w1&peJlRNlUF*-~(=2y47$Yzk#
zz@WgvHbr9(qv(Ny8BPo#hhSnMnkNp5o&kx8JcfyhXkBR(y8#lr!puH-2LC$FgIXsR
zG?_JUB~A7h(6DzFPdLKB=peweqWnQ?07q4#i6Vp6NmdC%2L=HaW-qzAMiHh0315ze
zq%uY~9svfHgAJ?L3oTj`WZZ<A7R;NxP2d3I!O7)<2U(8#y>OXqBlLoqQF{Gkbzyy>
zB4(a6$sR=sLj#~u&8&7Q%9}HUe=*kgbX^f}U|qRLT`{_8^??pY<^~0ch)0jKDuj-7
z321K=(%w`gyvcO~SC51V14mP%tKNo1>Ic_wY&O!~?4<3k80`kv^h`;+M(9dcfcDlP
z?QKP(n_73kG_~q&d8B@DEywmE?d?t4;)>Cy;F|1TC221ZdeSvPd*>wWT}6VMj!uB-
nI;yurN$cP`j@_HIcOTL=V}Hri{IJa-{6NCA#z-a>1_%HEy{xqK

delta 550
zcmZ3Wvp`40CBn&h$_Cpx2Oq8~`1i#6nb!~1I}=4Z8JRZDD`w>2U}O_)Qe+g|EXlN!
zQC>;pfx|=*g#w3zA{h=7#YIva8ZXRXRIp`ss9|e3#54ILb0g#I&BZKdnFW~Bq9!H3
z$nPqapMA2x<9^8G0FG}QP7KW)2PZUeZ1&?k!pJ1hxLJfdnOU%lku|VLHITD4P@y%*
zfK@7BatGfkMuW-5{OWcG*$k2z7!)|zrfBS86g_Y-!-*ke6HF{b^Ta{XGaxaM%P=t!
ztt*XUH$Y-nKC?}p!M~2PQR~ElCbI^vgvtH_8urTK2}c+h9Rzq*ls{+<;HXM8QDo3M
z$SPszz#zcF>?K#%D8f`A;mgsGRKn=SBf!9NuwfN@p+#$gjGHji0=dcC1P(AZPA(Tb
z$kOW9;5ykx=mqmfsjZXMh4qDUn0eA9dlV%M4S+^9v)Uy+-kc%)i?P0=>xzH_>xxC{
ziqTH14|F&(Hz-I%Jo=<nA#|imKzpN*_NF4?O|EOWdL&F3IGP$=^)@b2Ke&ctvyt{@
zCvA7dXg0W}XGYpJLRY#1w6_LnZz~er)Vc+xsa0?5BlUx8Ikp#RZ*S5TSB&n0YqEb;
zq`g4sN!JAJos+b86$x%SIsvBZsNPN`t%K`0c5l+&eMsAk{UuZL!#0KR0}0a_Bbitj
GAOHXZ@V#OH

diff --git a/piet-gpu/shader/gen/transform_leaf.hlsl b/piet-gpu/shader/gen/transform_leaf.hlsl
index 7744e0f..38136c9 100644
--- a/piet-gpu/shader/gen/transform_leaf.hlsl
+++ b/piet-gpu/shader/gen/transform_leaf.hlsl
@@ -47,14 +47,14 @@ struct Config
     uint pathseg_offset;
 };
 
-static const uint3 gl_WorkGroupSize = uint3(512u, 1u, 1u);
+static const uint3 gl_WorkGroupSize = uint3(256u, 1u, 1u);
 
 static const Transform _224 = { float4(1.0f, 0.0f, 0.0f, 1.0f), 0.0f.xx };
 
 RWByteAddressBuffer _71 : register(u0, space0);
 ByteAddressBuffer _96 : register(t2, space0);
 ByteAddressBuffer _278 : register(t1, space0);
-ByteAddressBuffer _377 : register(t3, space0);
+ByteAddressBuffer _376 : register(t3, space0);
 
 static uint3 gl_WorkGroupID;
 static uint3 gl_LocalInvocationID;
@@ -66,7 +66,7 @@ struct SPIRV_Cross_Input
     uint3 gl_GlobalInvocationID : SV_DispatchThreadID;
 };
 
-groupshared Transform sh_scratch[512];
+groupshared Transform sh_scratch[256];
 
 Transform Transform_read(TransformRef ref)
 {
@@ -167,7 +167,7 @@ void comp_main()
         local[i] = agg;
     }
     sh_scratch[gl_LocalInvocationID.x] = agg;
-    for (uint i_1 = 0u; i_1 < 9u; i_1++)
+    for (uint i_1 = 0u; i_1 < 8u; i_1++)
     {
         GroupMemoryBarrierWithGroupSync();
         if (gl_LocalInvocationID.x >= (1u << i_1))
@@ -184,11 +184,11 @@ void comp_main()
     Transform row = monoid_identity();
     if (gl_WorkGroupID.x > 0u)
     {
-        Transform _383;
-        _383.mat = asfloat(_377.Load4((gl_WorkGroupID.x - 1u) * 32 + 0));
-        _383.translate = asfloat(_377.Load2((gl_WorkGroupID.x - 1u) * 32 + 16));
-        row.mat = _383.mat;
-        row.translate = _383.translate;
+        Transform _382;
+        _382.mat = asfloat(_376.Load4((gl_WorkGroupID.x - 1u) * 32 + 0));
+        _382.translate = asfloat(_376.Load2((gl_WorkGroupID.x - 1u) * 32 + 16));
+        row.mat = _382.mat;
+        row.translate = _382.translate;
     }
     if (gl_LocalInvocationID.x > 0u)
     {
@@ -202,20 +202,20 @@ void comp_main()
         Transform param_10 = row;
         Transform param_11 = local[i_2];
         Transform m = combine_monoid(param_10, param_11);
-        TransformSeg _423 = { m.mat, m.translate };
-        TransformSeg transform = _423;
-        TransformSegRef _433 = { _278.Load(36) + ((ix + i_2) * 24u) };
-        TransformSegRef trans_ref = _433;
-        Alloc _437;
-        _437.offset = _278.Load(36);
-        param_12.offset = _437.offset;
+        TransformSeg _422 = { m.mat, m.translate };
+        TransformSeg transform = _422;
+        TransformSegRef _432 = { _278.Load(36) + ((ix + i_2) * 24u) };
+        TransformSegRef trans_ref = _432;
+        Alloc _436;
+        _436.offset = _278.Load(36);
+        param_12.offset = _436.offset;
         TransformSegRef param_13 = trans_ref;
         TransformSeg param_14 = transform;
         TransformSeg_write(param_12, param_13, param_14);
     }
 }
 
-[numthreads(512, 1, 1)]
+[numthreads(256, 1, 1)]
 void main(SPIRV_Cross_Input stage_input)
 {
     gl_WorkGroupID = stage_input.gl_WorkGroupID;
diff --git a/piet-gpu/shader/gen/transform_leaf.msl b/piet-gpu/shader/gen/transform_leaf.msl
index 9c7e6b7..6a55784 100644
--- a/piet-gpu/shader/gen/transform_leaf.msl
+++ b/piet-gpu/shader/gen/transform_leaf.msl
@@ -127,7 +127,7 @@ struct ParentBuf
     Transform_1 parent[1];
 };
 
-constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(512u, 1u, 1u);
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(256u, 1u, 1u);
 
 static inline __attribute__((always_inline))
 Transform Transform_read(thread const TransformRef& ref, const device SceneBuf& v_96)
@@ -214,9 +214,9 @@ void TransformSeg_write(thread const Alloc& a, thread const TransformSegRef& ref
     write_mem(param_15, param_16, param_17, v_71);
 }
 
-kernel void main0(device Memory& v_71 [[buffer(0)]], const device ConfigBuf& _278 [[buffer(1)]], const device SceneBuf& v_96 [[buffer(2)]], const device ParentBuf& _377 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
+kernel void main0(device Memory& v_71 [[buffer(0)]], const device ConfigBuf& _278 [[buffer(1)]], const device SceneBuf& v_96 [[buffer(2)]], const device ParentBuf& _376 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
 {
-    threadgroup Transform sh_scratch[512];
+    threadgroup Transform sh_scratch[256];
     uint ix = gl_GlobalInvocationID.x * 8u;
     TransformRef ref = TransformRef{ _278.conf.trans_offset + (ix * 24u) };
     TransformRef param = ref;
@@ -234,7 +234,7 @@ kernel void main0(device Memory& v_71 [[buffer(0)]], const device ConfigBuf& _27
         local[i] = agg;
     }
     sh_scratch[gl_LocalInvocationID.x] = agg;
-    for (uint i_1 = 0u; i_1 < 9u; i_1++)
+    for (uint i_1 = 0u; i_1 < 8u; i_1++)
     {
         threadgroup_barrier(mem_flags::mem_threadgroup);
         if (gl_LocalInvocationID.x >= (1u << i_1))
@@ -251,9 +251,9 @@ kernel void main0(device Memory& v_71 [[buffer(0)]], const device ConfigBuf& _27
     Transform row = monoid_identity();
     if (gl_WorkGroupID.x > 0u)
     {
-        uint _380 = gl_WorkGroupID.x - 1u;
-        row.mat = _377.parent[_380].mat;
-        row.translate = _377.parent[_380].translate;
+        uint _379 = gl_WorkGroupID.x - 1u;
+        row.mat = _376.parent[_379].mat;
+        row.translate = _376.parent[_379].translate;
     }
     if (gl_LocalInvocationID.x > 0u)
     {
diff --git a/piet-gpu/shader/gen/transform_leaf.spv b/piet-gpu/shader/gen/transform_leaf.spv
index e561e9dcc5b5efa2b6e077e953d952b180e52c0a..f418bbef9311d0422c800bbcce4d1f3d78e2a24b 100644
GIT binary patch
delta 2823
zcmZ9NYj0Fl6o&V7+FFVTZKFXOp@^u7mlzXCh=Byu6cQsXifFymskRh*b(jh&b-1Z`
zL+T)kh=@15AOaN;75w5KAbvA`@CW$K#7h$spXbck@svqd*89F|?X~w=d!5tY3%?d-
z<!csR6hbaEhKBIt#H`iL!?_UVg!)h$LTJK5?!@r28EYRv9x3(r7E$U57T}|WVqqxd
zk4@wGFbFS>Y$CTlSaSVf_qvhdKzngyY;@Hf)~E}mO0KbGX^E)6p+bq=`XE;-6$`@~
z2MeWMu-Q)`b`*-e!=)8teQ5P+mCnYN^5(RH=XwR4J=bf~cs@+PHzXDFVLv#Y|G>1o
z+Hi0huM3Ca$%*<#iq$oB;Yd}xX?SCIvDYJcfD_Xe)`rs;@H3N5Gk>lr+mo3rxdL0s
z!q|m{aDP!FEvU`G+9^MVP!DdyR>JLyzqy!eB9C?2Lq5^+#9eQ%uBvYPzp#K{OgJ|=
ze#w^FT&^<J+*R%*y8*k6m5q2$&0m-Y_r&wS^RX*g?#ZfJbw_y{d7JPx<ajdc$UTLc
z@hvHK?e*}L=%0dTcuEsarU}jX8+iB`_;ZziTmC3JluU0@_7=QDndJ^;mOGSL?yzLJ
z!xA~$C&LoC^&OZD?+H##Mo6Y6^6OaO)I@#_+=+?20q(FwZrovs-1TQB&tG~d-_(J;
z8k>zx(UFeIw#(<w@^Y4Q<rm}ERL;+B8M;S=EhWbxxfi?&`-v?1T5x}=A46*d--qTp
zV}Egt_v8J!kFj-_+TUE5IPTcxj^zhYy0Ez8L*VR=55wJ2eS3O}N5HN#X0>*FH?{&Z
zX00~XJPKA{#;Lq%d5rad(aK5CYgdTD$HDV37dis+B{aaIz2L=o*BCR$T=jo=#Q3QC
z-#ExHxmvc_eZ+3ZdP3!wxsA832P(zw`@xos&$a>Vvzh01^8U(rYwMyXKz3_u+L-~c
z+8#DE)7U`eZ0r2wMO>2hWeD}=G{cFO$6FZz`;zt$_dMSEC|vz$dgD)m)y(n6H-i0n
z<Hm}Znl~<{b~zccry&;>e+ErGUhcDCH7|WR+xvEpW1F$N(TsaA-}dJ)W1Sq#FG)Go
zfXM^x#&1O%%ym~6+Y<#7UF4@{Te=9BEZT?P4zV+_k8<<8;P_<v*tH$b@~WLEJFwWN
z=fP^RPcMS~*(YN!U~2YB9INe0-i((}c44tGFN4)=Oq`Tg;hxjy@>?XgN;`fxwgPiq
z%zr)A&G%a+kAuAj>|k4mSoO_Fnt<Sg$D`^kFoE~cjJ*N&G4y}(czf}0V*ZTpOXK=D
zR&Rl=aXcNXx4~+&v8hn54GI;^Vip(2#qWS!9G~o6aO}hZu-s1g)qM|qHRhU_e+aDZ
zgB--mV<+ARFD|ptLIzFRk2hdZ=jsEnnloup=kG(fx;c)~M`?~@WRBzZF`9Y@IoHBJ
z0lQA!eGcQ*%FFqX-$lY_kjpSP^50tA&^PmQxN(cizW_I6_))N3ir<GX!RDy%<N&TY
z23B7Kj(J~s{}wmcLGeWdkAn#$r#ym}`-YyxPh#0adZ1>suhW<>o12|VHDJ;u|4Dv>
zb^`Ok?121Kj{aN3jlV^3LxZRBa$~mdJFqP?b_OqBLg#9TeGk47Z*kWdcNiT^=Sd$s
c_X9ZFsk07@V*CUcJd0(geedL*3y#eG54^fawEzGB

delta 2817
zcmZ9N+iz4=6vp>-+S(QoXrn<I3y7%k7GnYlF_08Z3W<>xK@^la)s|LzfoUtE;2`1!
zF9?jFB7))tP(YxHST9K8ldrz|;2)q*CSG{*_d9cTJY~|A^?l!3d+oi~S?BcE!taI2
z`Kmcrgpdmjp)UNmXY#Vfp<D=4Lv5%IAxy(U)tSLX^($8+4~_J87g1^mcEL9kiiLrc
zKQV#l!vMTEw29o>V99m;ovVh5>)VP$qZ^jqYmJ&PTFx~zEgU84Z=f(jZf%e&jT8%m
z!~KPkZm`*#5ZepI?!l48qdjP~YUPfGrqbrLg6Dc2oITgB2|OS6!uyhn`EURn&wp@2
zUUfJ$f!Bm1@Z>~2L&eIPnsBV5-7z@aS?u;m9^mYRh1KEwWqf>W+N58qO7>(DOJ-n8
zSQxu72ktLwqy@F9SR3V^LZ}7L!j{18iofZYYa)+z+Co0j^2A+ludc6b`@b-oU`)6$
zw)Ltl)w!zjrN-5zcCzaU-pWdMT@J6q?uIwuy#;@v1?~-6z%#MyDCbElTIB@tS>!#9
zuOi2jT1W1w)QoRUF>7y0HeN#DIHCtJVJuB(#NWoT>hTxKe>DABa!8rxr0iXIhcwF_
z(kypKv)tjya)&2!cXxOqx4r|E;a$Ng$_U96MSe32oTA8YggZfz*TEg0$c;NZk-L6;
z?Bdmj^V2$zmtj7hOLU~8yz$zZlYKa?jKG!n2g(<xHw~;1VGGIe@vQ_e#eO17{vfzF
z)lZ={fLEcp&e-2P#)t9#+{f5MnA$%)Fmc>*jXRbeL0N;v9Ulc}cU%j1NA>OLDLTQf
zGiJ3m{A1W+%$T)iv8Di4U&N_=(()L49E?^<f?W`W80-eOU@mk7<O^tkMSH;W@vbpu
zj=Adpa>aO6{cjv(kX$X>>~+NM!MZ~E=jjdedVxxD`#!KG<K3PFdpGmkPTp4@Xl|an
z9%Q$+rJd;qtL<h(GmZ6^Pd3jix&4B)FN5Ga(+nqC9-qnvurFyBanIvZ-w0Pfkv{Pv
zSj`-t_z2jaPu$orrsfkDQ@fN5*)x#Ki$9B|9v|)&u$m9OmF<1IpTjm|_oEs2V16k|
zn6VBH=9gq!ssWP=?ZQ8gHjwMAEVd^KCc4N^&-Qc?E?Kl6{{qC$#6HT+^8v@3>0{S+
zILk|4MA?DGKD`81i+y?p?9V<Kdl^%+PvTf@XYyvein0@njd=~MW@F-{yaD%|-pg;1
z+$wGOH?hT->tg=yR5#yml{^mi9<YOL9b%O?!{P?Kcs#1!0(<dZ%~%=iHS~XRy?yxo
zm_Orhr*VB8t9QWGIGv8wyI{2`*riaa4hnB!7PGiGE`AT};&`+7!Lbtu!E!s{SNAaZ
z2Fx`v{{yhP7deEN$4-0*o?l|2g$$Z>0B^ve&eca?HD}VI&fmvyb#ok}PtqL6$Q;M*
zQ#AE<a&Cej1-nk&eU9MOO09gz?<3(e$TgT79Z$W*d^10X8;|@1xGuv_g6&fLKAZxZ
zqrQU!SmO(@`f_m0JMHtgxWRTdn@jKvm_TyMV|cl5=otP>EPF@~)QI*~8uMjyvva8i
zOuFPh$*<AQVqVY=$j{|gRBq+Q-ypc5!Si^z8`!>Y!M4oUI9|Sh&Q%fn4tzV_;;tLd
fmDje+nD;#}+ms*RYVl<^rk0)SePi=xAD!|a=;KGS

diff --git a/piet-gpu/shader/gen/transform_reduce.dxil b/piet-gpu/shader/gen/transform_reduce.dxil
index 1ed5e0e3dc4b5a533f1829c1818e590d6ddede9d..63df381194f6336e5c7fa13118ab5248ff4ef02a 100644
GIT binary patch
delta 270
zcmV+p0rCFWB-kVrL|8&YWXuVB(4rr|An%((ZBwOJkrasm0kNK20SJfy1`&cG0SU7K
z0<{4I5)>d<le+^MlhXr+0S2>L1d0O?(gC5`7nLH^NtY{6YMEb6lk^An2qFN32*Hqm
z1heu8zySh~fU~p;QUeV^lw5&SLSYn2fd-UBlM@ah9>Fw(K{o~?1_=OyU=X1JRgA$H
zkim#S0)QYGL}^fyQVuZznv;YMB^5CSFk*lJAP5GO9h89+n89d508}EAz78h=bCc!{
zS^<KSK@TYwAO#pPKmZT~1L_t`;SGXQ7)=O(N=B1_4=4f$6_ctDb^*bYArQR;1W5ct
Ulj0ES0(C)?&JiyGfRp|aTAeRg82|tP

delta 288
zcmV+*0pI@EB-kVrL|8&YenxVBcy;w0`_d_NYJZekkrasm0<oT30SE{I1`&cG0THtS
z0<{4eA`~ESK#>$wSq_v|X_Nwu8358V0K{hiD9n@41Bn49vswg-0}u$^i;6p+PA7CH
z?r^3~fB}>A2lfa;0D}m@kbnrY@(92I0uX_-v<gxK4apTmfs{lg6h@&GXh4$_4j~=^
zHH1Mo1|tRu0D@o;p+QxQ!5EOih(Q8?AQ(hxV3Se~F#(yAgbpPYAqFsFfB+x}29zC?
zffJa)XhHx~5|h3TCjn!V<_=l`0+T@xDHa$77%@Nq5CjA27EIv{f>Rhx2!KjNlYkE>
m0_PNyst<Mnfs-K+y#&NRj6{>-5b6SSfs@V=F99Hv{t;Sc99<y*

diff --git a/piet-gpu/shader/gen/transform_reduce.hlsl b/piet-gpu/shader/gen/transform_reduce.hlsl
index 5ada811..af52cdb 100644
--- a/piet-gpu/shader/gen/transform_reduce.hlsl
+++ b/piet-gpu/shader/gen/transform_reduce.hlsl
@@ -36,12 +36,12 @@ struct Config
     uint pathseg_offset;
 };
 
-static const uint3 gl_WorkGroupSize = uint3(512u, 1u, 1u);
+static const uint3 gl_WorkGroupSize = uint3(256u, 1u, 1u);
 
 ByteAddressBuffer _49 : register(t2, space0);
 ByteAddressBuffer _161 : register(t1, space0);
-RWByteAddressBuffer _251 : register(u3, space0);
-RWByteAddressBuffer _267 : register(u0, space0);
+RWByteAddressBuffer _250 : register(u3, space0);
+RWByteAddressBuffer _266 : register(u0, space0);
 
 static uint3 gl_WorkGroupID;
 static uint3 gl_LocalInvocationID;
@@ -53,7 +53,7 @@ struct SPIRV_Cross_Input
     uint3 gl_GlobalInvocationID : SV_DispatchThreadID;
 };
 
-groupshared Transform sh_scratch[512];
+groupshared Transform sh_scratch[256];
 
 Transform Transform_read(TransformRef ref)
 {
@@ -101,10 +101,10 @@ void comp_main()
         agg = combine_monoid(param_4, param_5);
     }
     sh_scratch[gl_LocalInvocationID.x] = agg;
-    for (uint i_1 = 0u; i_1 < 9u; i_1++)
+    for (uint i_1 = 0u; i_1 < 8u; i_1++)
     {
         GroupMemoryBarrierWithGroupSync();
-        if ((gl_LocalInvocationID.x + (1u << i_1)) < 512u)
+        if ((gl_LocalInvocationID.x + (1u << i_1)) < 256u)
         {
             Transform other = sh_scratch[gl_LocalInvocationID.x + (1u << i_1)];
             Transform param_6 = agg;
@@ -116,12 +116,12 @@ void comp_main()
     }
     if (gl_LocalInvocationID.x == 0u)
     {
-        _251.Store4(gl_WorkGroupID.x * 32 + 0, asuint(agg.mat));
-        _251.Store2(gl_WorkGroupID.x * 32 + 16, asuint(agg.translate));
+        _250.Store4(gl_WorkGroupID.x * 32 + 0, asuint(agg.mat));
+        _250.Store2(gl_WorkGroupID.x * 32 + 16, asuint(agg.translate));
     }
 }
 
-[numthreads(512, 1, 1)]
+[numthreads(256, 1, 1)]
 void main(SPIRV_Cross_Input stage_input)
 {
     gl_WorkGroupID = stage_input.gl_WorkGroupID;
diff --git a/piet-gpu/shader/gen/transform_reduce.msl b/piet-gpu/shader/gen/transform_reduce.msl
index ac586d9..c387f03 100644
--- a/piet-gpu/shader/gen/transform_reduce.msl
+++ b/piet-gpu/shader/gen/transform_reduce.msl
@@ -72,7 +72,7 @@ struct Memory
     uint memory[1];
 };
 
-constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(512u, 1u, 1u);
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(256u, 1u, 1u);
 
 static inline __attribute__((always_inline))
 Transform Transform_read(thread const TransformRef& ref, const device SceneBuf& v_49)
@@ -105,9 +105,9 @@ Transform combine_monoid(thread const Transform& a, thread const Transform& b)
     return c;
 }
 
-kernel void main0(const device ConfigBuf& _161 [[buffer(1)]], const device SceneBuf& v_49 [[buffer(2)]], device OutBuf& _251 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
+kernel void main0(const device ConfigBuf& _161 [[buffer(1)]], const device SceneBuf& v_49 [[buffer(2)]], device OutBuf& _250 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
 {
-    threadgroup Transform sh_scratch[512];
+    threadgroup Transform sh_scratch[256];
     uint ix = gl_GlobalInvocationID.x * 8u;
     TransformRef ref = TransformRef{ _161.conf.trans_offset + (ix * 24u) };
     TransformRef param = ref;
@@ -122,10 +122,10 @@ kernel void main0(const device ConfigBuf& _161 [[buffer(1)]], const device Scene
         agg = combine_monoid(param_4, param_5);
     }
     sh_scratch[gl_LocalInvocationID.x] = agg;
-    for (uint i_1 = 0u; i_1 < 9u; i_1++)
+    for (uint i_1 = 0u; i_1 < 8u; i_1++)
     {
         threadgroup_barrier(mem_flags::mem_threadgroup);
-        if ((gl_LocalInvocationID.x + (1u << i_1)) < 512u)
+        if ((gl_LocalInvocationID.x + (1u << i_1)) < 256u)
         {
             Transform other = sh_scratch[gl_LocalInvocationID.x + (1u << i_1)];
             Transform param_6 = agg;
@@ -137,8 +137,8 @@ kernel void main0(const device ConfigBuf& _161 [[buffer(1)]], const device Scene
     }
     if (gl_LocalInvocationID.x == 0u)
     {
-        _251.outbuf[gl_WorkGroupID.x].mat = agg.mat;
-        _251.outbuf[gl_WorkGroupID.x].translate = agg.translate;
+        _250.outbuf[gl_WorkGroupID.x].mat = agg.mat;
+        _250.outbuf[gl_WorkGroupID.x].translate = agg.translate;
     }
 }
 
diff --git a/piet-gpu/shader/gen/transform_reduce.spv b/piet-gpu/shader/gen/transform_reduce.spv
index 5638afb4fd09f1c3370af110cf9954c57d0ba53e..af5ffb9175997a0b9a18f45cf6ddb417f3e1cccc 100644
GIT binary patch
delta 1621
zcmZ9MTT4_?6vy|OnKO=KmO7MBG-=XX5#dvY-3Y7)ksuN5K@_y0#*x#^gD*`z^a0`&
zLEj*s68!>^)k73TcF(jdvwL>&*7g58b9Pvp`LovgughL*@3Uv5ak#NO8mTKW#xY5g
zFw0Jtw^V&M##ETNiAh!p#u@A#a9f{B%jOPudy=l4Q<8n1US}q>r~k2CGb{Um*V%il
zC+iVIK{!^XGbcN4X8n`S%7Za8FRmq@+nw(r&JAqhviV#WoNJ<H5kkvV1T7vq+R>Kv
zjx>4MeBZu3#hwz5G0okXtk+MSxJme9@#;;m`k733rlZ^QvP{4(@R1@PWUxzesJ5gn
z;>`MkC8-UxOJbc+Av9@(Imx(CBkZMBf3P&w%m7h=2pU~4u&iAx-TG<1OjsTIOSA<Q
zId3W^{K{mi)i%xEe?5w%VYr0Qq?*Lp%1P-<cgC(&9E(oMhWeASsN?v@%i9ZgMUtGB
zgFzXP>t3sJ3<;N(&S=;*id&GR0l00U&A4#bXC)a=WrO(NlK7}aenbvz?bRL-Zb;aT
z?OCM5>)~7gJOxU_cp7IzY^J2>Mn#^HEb}idpRDC`)ps34RZC9&XHO4ZS+O}<IV$>&
za9&tc!BO8$tt;njcBzGPlK1>AsZ`-Un$V~i8hRky5JnV*{ZMES;7^3yBkAPfyC`m0
z^05GaS(wx5j!ClE8{ZQFZeAmyhh0R>c+kd#sByuz@l5*CHlE9-4S3ter1S{^ALXt~
zz7Vbo_^4GS`BH%2uC_QmZ0ma^{em%tz{ycUSm$f$B#_9z0QOd$3Ps-tYXoxe5u-u@
zzn~fItjGy9vQrW)y@9u~uL<?WZ%8Ez@1#Q6?cYnMq@C@9bY>%t?XZbqyW5f+Et34W
z<VS%swCD9H<hmru#|KCLO{1wOa7%>!LS}-i`7CBe7zl6ri*&epft&s&$>OH*eHGxS
zhHk6<kj@(-&$jqeI=mhEm-OPual;$agDV98J89_z?^hFgWX93kAehAl5!q2e33}Yc
l`LKy2hj$p0PL4#7!yAob2=+{{iQ`O)Lqw$vom#ac@dvyh=U4y$

delta 1637
zcmY+ENl#N@6o$Xf15|-VNi<d}apMpdtkmEX6Ldj{iG)NK1d^grOUack*jW7u=4$*6
zvXT^kfkYQZqfwku6epa;p@8H2+}_XBo4z^cyl4Ktdpq7f(VhwgTVuuqOwz<n>U65B
z>8CNK!9-11vQaR>vBA+u_Y-M_;lnvcG9u?I$)P@{FYnv4|Je~!k-gvP8$3QxaEPHG
z98bP)Sa#G@+*5&N2g7DTTvu^;S8;&2h-VX5C=U0-MNG*2fY1|B1S1|g*4tBXj&?YO
z;?UmR)tM3jV>)yBf-^#$s7bg}(dLb?`k8z#-<xxs0vpf;K3L_w47${nv$39FpyH0j
z(rX!)#8#m}=+FxDl2M^W*u$vqSR&oY0wIA223;fYSUW4-`We1XSmFEQYC%QLn~Djy
zF`4eRL(}`8N0AH+7Z*BIlQ>(sPI_%;?5yH=n54cuk_?9eLH9)JK<T!~gs@%?+TRfv
zB_o@~S(xizrSdEdmyyn@*e!}%lw=^dt-j6raM))hSyyGfcocl9b;Kh=K0z}IMt$2h
zwr7*^uZuweJn>0Co>QP9He1tVgCfsJ*1Ka%X0p7#de6P6W=US(T6?y9ap|T|<D}?Y
z0;jX6f|KsIrqwCVW~W9tCwbS+rqiW+XhJ(fOWqf*2@{IKejscS;7|J8L+Rw<yC7~{
z@{s_4NvP;_A4~Gk8{d=wH?Nh@{V9TG+8g7ksA<8D@l1N{7&EdN1Ky7DTsnF9D0fxz
zg>YHGN3ABwmje7Yjm7C<Ti<Kx=Zz_OPL2}7I^Reqfkgflus7>eDEd}dDUgGY7!~sP
z1?^~eMNVsw9g?8-2HwfOBD5KIT{>BMFBQtR{~(=`cDIkx*^M~bVG~2Uo01$YlKiOT
zCxJ7x=k?j=`X$N72S@&Oqpc`#Lxg=p=De%<B4$n)^>6yCbhtKwoBl4z!%gG+Ccse*
z-B$Z0oo|RdJK}HY@OI@t(yJ?L0lof#fUbrQF(VMXPebsH1?41|hX#Cq;c9{x%YN9z
jk<+NSh;(u!yd1vKn4l)co(VQ_oJn<wkd$(9`S$o<N5JRN

diff --git a/piet-gpu/shader/gen/transform_root.dxil b/piet-gpu/shader/gen/transform_root.dxil
index 0d16d04a7a85d55f33415f8b69781288a088677c..5b4f05930600267411a9865e166d35bb5c6578ff 100644
GIT binary patch
delta 302
zcmV+}0nz^0CD<htL|8&Ydl$hAPL=UVXx-@7xYAU(krasm0kNJd0SJfy1`&cG0SU9w
z0geG35)|M-!4x1^8I(fd6*WPD6+uDJm;oR?13+O00JD4p_yZ8;UsfaQ9<+VF>ys=s
zb8)hhga`HrA^?L3!H|Fivw{er0RoVKvn2`_0}2IEhyevspaCUOlSU0!C)bS_AOHx0
z0ksB~5^SRh0Wg{%1fvOqFq%LNWo$cQfB+x}22>n$lf(_KD}_C9gWwNF69Ql~K?p_@
z24OUT7>)q;2w;x@_6T508HF7&KmZT~11c9vlbsHz0p^oN5A*|}5-F2C5b6ZRU--q7
zD-kXe7!5F(WiW<d#J~VR5Db>w2o2D)Y!N^K0kD&;5*-d+1Po>wlq5g^0000000hxt
Aa{vGU

delta 302
zcmV+}0nz^0CD<htL|8&Y#>My)W%}jMucRe0y94~ikrasm0<oSe0SE{I1`&cG0THv&
z0geG3A{5|2!4x2HKpB)m;T1JOffYeP(3k-rJp({t2D5wv_yZ6Pq~2NH0UVkcOT9BP
zu&Z>Fga`HrLI8sZ!H|Fmvw{er0Rj+#vn2`_0}6#0P!a`EAO#vwlSU0!C)JG@AOHx0
z0ksC05^SRh0Wg{%1fvOqFq%LNWoSEMfB+x}22>n)lf(_KD+NAqgWwNF69Ql~K?p_@
z24OUT7>)q;2w;x@_6T505rrKwKmZT~11c9xlbsHz0pycL5A*|(5-gKF5b6YmUxdb!
zD-kXe5DhSxWiW<d#J~VR5Db>w0FBYJY!N^K0ictu5*-d&1Po>wlq5g^000000L(gH
A8UO$Q

diff --git a/piet-gpu/shader/gen/transform_root.hlsl b/piet-gpu/shader/gen/transform_root.hlsl
index 35961b1..d447db6 100644
--- a/piet-gpu/shader/gen/transform_root.hlsl
+++ b/piet-gpu/shader/gen/transform_root.hlsl
@@ -4,7 +4,7 @@ struct Transform
     float2 translate;
 };
 
-static const uint3 gl_WorkGroupSize = uint3(512u, 1u, 1u);
+static const uint3 gl_WorkGroupSize = uint3(256u, 1u, 1u);
 
 static const Transform _23 = { float4(1.0f, 0.0f, 0.0f, 1.0f), 0.0f.xx };
 
@@ -18,7 +18,7 @@ struct SPIRV_Cross_Input
     uint3 gl_GlobalInvocationID : SV_DispatchThreadID;
 };
 
-groupshared Transform sh_scratch[512];
+groupshared Transform sh_scratch[256];
 
 Transform combine_monoid(Transform a, Transform b)
 {
@@ -55,7 +55,7 @@ void comp_main()
     }
     Transform agg = local[7];
     sh_scratch[gl_LocalInvocationID.x] = agg;
-    for (uint i_1 = 0u; i_1 < 9u; i_1++)
+    for (uint i_1 = 0u; i_1 < 8u; i_1++)
     {
         GroupMemoryBarrierWithGroupSync();
         if (gl_LocalInvocationID.x >= (1u << i_1))
@@ -79,13 +79,13 @@ void comp_main()
         Transform param_4 = row;
         Transform param_5 = local[i_2];
         Transform m = combine_monoid(param_4, param_5);
-        uint _209 = ix + i_2;
-        _89.Store4(_209 * 32 + 0, asuint(m.mat));
-        _89.Store2(_209 * 32 + 16, asuint(m.translate));
+        uint _208 = ix + i_2;
+        _89.Store4(_208 * 32 + 0, asuint(m.mat));
+        _89.Store2(_208 * 32 + 16, asuint(m.translate));
     }
 }
 
-[numthreads(512, 1, 1)]
+[numthreads(256, 1, 1)]
 void main(SPIRV_Cross_Input stage_input)
 {
     gl_LocalInvocationID = stage_input.gl_LocalInvocationID;
diff --git a/piet-gpu/shader/gen/transform_root.msl b/piet-gpu/shader/gen/transform_root.msl
index 2c58c06..8b4b2a1 100644
--- a/piet-gpu/shader/gen/transform_root.msl
+++ b/piet-gpu/shader/gen/transform_root.msl
@@ -62,7 +62,7 @@ struct DataBuf
     Transform_1 data[1];
 };
 
-constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(512u, 1u, 1u);
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(256u, 1u, 1u);
 
 static inline __attribute__((always_inline))
 Transform combine_monoid(thread const Transform& a, thread const Transform& b)
@@ -81,7 +81,7 @@ Transform monoid_identity()
 
 kernel void main0(device DataBuf& _89 [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
 {
-    threadgroup Transform sh_scratch[512];
+    threadgroup Transform sh_scratch[256];
     uint ix = gl_GlobalInvocationID.x * 8u;
     spvUnsafeArray<Transform, 8> local;
     local[0].mat = _89.data[ix].mat;
@@ -97,7 +97,7 @@ kernel void main0(device DataBuf& _89 [[buffer(0)]], uint3 gl_GlobalInvocationID
     }
     Transform agg = local[7];
     sh_scratch[gl_LocalInvocationID.x] = agg;
-    for (uint i_1 = 0u; i_1 < 9u; i_1++)
+    for (uint i_1 = 0u; i_1 < 8u; i_1++)
     {
         threadgroup_barrier(mem_flags::mem_threadgroup);
         if (gl_LocalInvocationID.x >= (1u << i_1))
@@ -121,9 +121,9 @@ kernel void main0(device DataBuf& _89 [[buffer(0)]], uint3 gl_GlobalInvocationID
         Transform param_4 = row;
         Transform param_5 = local[i_2];
         Transform m = combine_monoid(param_4, param_5);
-        uint _209 = ix + i_2;
-        _89.data[_209].mat = m.mat;
-        _89.data[_209].translate = m.translate;
+        uint _208 = ix + i_2;
+        _89.data[_208].mat = m.mat;
+        _89.data[_208].translate = m.translate;
     }
 }
 
diff --git a/piet-gpu/shader/gen/transform_root.spv b/piet-gpu/shader/gen/transform_root.spv
index 7824d0914594eb18dfe4cd7dac5988c2b784e2d9..1578842ea8970357fca4c1f6c7d85bd116242059 100644
GIT binary patch
delta 1715
zcmZ{k*-KPm7{<SuQ{!Y#Bwci27eRC(El|>hh=CWOMxjABEi{M@G~vc*MG%a1Q$a*>
zT5PwhY%y~hTWqlu+iW#kEDP#S=%x*Vp5Ju7I9)XG@IKGGeBb+?_nc|ZAIpz8!C6_x
z1WdLGn-8s#wd(??KX8m=g5ra6Oi)JYvS?{VMd`97(dx#=XjAi++E}zUULT9r#1paV
z=J=)<d0}5(yuP}2V~sB%mHZWm9oLx+&Ff<c*@>#?lx?g^RMmO<+&|hem!%U8o0aRB
ze(`uF&oTF9ZOub*{~5#X=U``GJiXJ|8Vm&7P-u0@ex)$Z99_9k_J9!9Xfvf>c`-eU
zjDG9su<TqDGD~I0-BhR}N%LeCsVX8%^L|b9eogaMmx(ig)zmYDUo(hb4|}>q<1aTR
zJ(~5!Dc&dNUZGO>sYcLygjV6ZhJ=5TpOFEhdIy*`H;^-}?w~wvLb0-V7>9)2!jP^(
z9~O!P{N3_xkUb($N9>z+aZHvU9S}Pz;G+Y$y|JAelPBb97wnBE#r-#SNOL3pZC)3r
z#Hl03X}e@k3)=-^oHa@IjDUZ!x?(|S8#^m5UosQyl)zzw=fztD3fUaGQ2XZ43qr0y
z4KZ?BWbyydh;~%${~mUiVmpKa<7Ve%FS;m-%U<6t-Xhr9E{QW6d0dC?c1I$4vwOsN
zS}op_Nr~f|FOW$r<>pSAzQDipLht#c6`ky5@<O8Rc6!-P`4xG3g)W`XuCEIC3Iuk2
zU6voaCU#B0$FAY_w7$%)Z^+Xp*gf79$H#u{thdCOm6&}Vx5e=@5q+Y0)b=JUv3^-P
zBX&oYF0fzK-GI^l1|^L9acW+2NQCqrZ?Jpf_$md4LEjZrmN^DaJA4EW1Zr)2kHl$<
z7WhuljKg>QSXe1gWAmTz+7X~&SXPd|uh&y?UM~(JMh<@l_-zBv#LIp9xi}5kx8Q|1
zIr#0RBjTJ>qQ|gC_nVgztArT>1uN8Xmh7lNGC!u^Kk&6cKDji6CXbr1>>GiafItoZ
qlVbuk_Ekia$8rYCMA!}su~WSh_fPR&8lOFY7(V}Y2h#I$o1EXAxBy-N

delta 1703
zcmZvcIcQW-6o&8QX)-a65JWAs5Uhe?p(sjFNT-5gz{ageA*c|PAXq3Q*k~b&j9Xk2
z<Gy4u#(m$SagF;5ww5+-AozXbyD*9me$Ickd(OG%&b+B=smdiKW6DCvgh3%2K5oe^
z8eXFI2T4c*YPgvRUA@0DOIqM7RxMr9fE*YGgrmrnbq#g(seb$~Jqg|D4J%d?n}lod
zWresT^de)<eR%O15As#Xk<5qg{$y=QrX)Y7bV2PlLdwBdN=`uT0JatNo$xF@kH-Hp
z^)XXHf0OE2<hW28W+6Z1KbN*Ow!n<4CM$<@m}1>wigkyH`V6=OMcw?~e{=^c_CHAV
z8m?n@2;CLgFG=Ngj9Y;t{-lxZ`P|?U!yHgJSQ;Lh^EYxs>vv&k2Gi23*bUzVZgZ$t
zvj<cI?d{l>BliOH^nGI!`;mUu)3*<3`_84xW5rf0HXgvz3Sz}JcyYyp=vLIeF74tF
z+&q1rwh?(4tOxo$%Y_~R+NaZ1OR7iTF?4K=1*f~w5S=IBO~6F=q&}7Hdgw_|3Cz)F
zOcPT34@QimDnHT4|1SF!z71ev$oJ(2O=^c~isd`tO(4$J33oQ*ET`U)FCQ|pz6<7U
zHKjLm8m?_TFjAj38{3R@2V0Q;Ha3+;9qtUM0OtB2)#JW;+}F&rSbD%oD!To1K-)y%
z_Aem)+`hi^K%152@wAJD?O(!j5yZ`2hHG=<ar!H8r`H#6=qg-$>{HztV<n}&Ye+lO
zcO7XL%CB^H)9v3v@iX8y(g3A}eDHUGhjgan-c4<Dk&dW-18`U&R~~ck!Oe|L_rbkX
z8@rF}qkdsZf8GzlJYZ5xc$C^r@aL@_pW$P;&(MSP8Pki@9vgT9pIM|oh1*bkE&bFF
z2I+`JpTRw-hWEY@`5ep#qd-5X<tECIFMzRruF8MLSHO5<ZAjfXbF#?Sz?=*)r_sO9
yx4@+MX4MUJC9cPnDC1P`;KftCN7okf^=T{Kx$F5w*(qQ$g?;P1^X@g3i<936GV&Jy

diff --git a/piet-gpu/shader/pathseg.comp b/piet-gpu/shader/pathseg.comp
index ec0a440..12104eb 100644
--- a/piet-gpu/shader/pathseg.comp
+++ b/piet-gpu/shader/pathseg.comp
@@ -10,7 +10,7 @@
 #include "pathtag.h"
 
 #define N_SEQ 4
-#define LG_WG_SIZE 9
+#define LG_WG_SIZE (7 + LG_WG_FACTOR)
 #define WG_SIZE (1 << LG_WG_SIZE)
 #define PARTITION_SIZE (WG_SIZE * N_SEQ)
 
diff --git a/piet-gpu/shader/pathtag_reduce.comp b/piet-gpu/shader/pathtag_reduce.comp
index 86bb9d0..9e84bf8 100644
--- a/piet-gpu/shader/pathtag_reduce.comp
+++ b/piet-gpu/shader/pathtag_reduce.comp
@@ -11,8 +11,8 @@
 
 // Note: the partition size is smaller than pathseg by a factor
 // of 4, as there are 4 tag bytes to a tag word.
-#define N_ROWS 4
-#define LG_WG_SIZE 7
+#define N_ROWS 2
+#define LG_WG_SIZE (6 + LG_WG_FACTOR)
 #define WG_SIZE (1 << LG_WG_SIZE)
 #define PARTITION_SIZE (WG_SIZE * N_ROWS)
 
diff --git a/piet-gpu/shader/pathtag_scan.comp b/piet-gpu/shader/pathtag_scan.comp
index c0d386e..7c1e74b 100644
--- a/piet-gpu/shader/pathtag_scan.comp
+++ b/piet-gpu/shader/pathtag_scan.comp
@@ -5,10 +5,11 @@
 #version 450
 #extension GL_GOOGLE_include_directive : enable
 
+#include "setup.h"
 #include "pathtag.h"
 
 #define N_ROWS 8
-#define LG_WG_SIZE 9
+#define LG_WG_SIZE (7 + LG_WG_FACTOR)
 #define WG_SIZE (1 << LG_WG_SIZE)
 #define PARTITION_SIZE (WG_SIZE * N_ROWS)
 
diff --git a/piet-gpu/shader/setup.h b/piet-gpu/shader/setup.h
index 5d4cc73..4211b08 100644
--- a/piet-gpu/shader/setup.h
+++ b/piet-gpu/shader/setup.h
@@ -27,6 +27,7 @@
 
 #define GRADIENT_WIDTH 512
 
+#ifdef ERR_MALLOC_FAILED
 struct Config {
     uint n_elements; // paths
     uint n_pathseg;
@@ -59,6 +60,7 @@ struct Config {
     // Offset (in bytes) of path segment stream in scene
     uint pathseg_offset;
 };
+#endif
 
 // Fill modes.
 #define MODE_NONZERO 0
diff --git a/piet-gpu/shader/tile_alloc.comp b/piet-gpu/shader/tile_alloc.comp
index 3761e9e..024f499 100644
--- a/piet-gpu/shader/tile_alloc.comp
+++ b/piet-gpu/shader/tile_alloc.comp
@@ -72,7 +72,7 @@ void main() {
     // Prefix sum of sh_tile_count
     for (uint i = 0; i < LG_TILE_ALLOC_WG; i++) {
         barrier();
-        if (th_ix >= (1 << i)) {
+        if (th_ix >= (1u << i)) {
             total_tile_count += sh_tile_count[th_ix - (1u << i)];
         }
         barrier();
diff --git a/piet-gpu/shader/transform_leaf.comp b/piet-gpu/shader/transform_leaf.comp
index e158c50..c51dfe6 100644
--- a/piet-gpu/shader/transform_leaf.comp
+++ b/piet-gpu/shader/transform_leaf.comp
@@ -10,7 +10,7 @@
 #include "setup.h"
 
 #define N_ROWS 8
-#define LG_WG_SIZE 9
+#define LG_WG_SIZE (7 + LG_WG_FACTOR)
 #define WG_SIZE (1 << LG_WG_SIZE)
 #define PARTITION_SIZE (WG_SIZE * N_ROWS)
 
diff --git a/piet-gpu/shader/transform_reduce.comp b/piet-gpu/shader/transform_reduce.comp
index 4b72b11..e59d559 100644
--- a/piet-gpu/shader/transform_reduce.comp
+++ b/piet-gpu/shader/transform_reduce.comp
@@ -9,7 +9,7 @@
 #include "setup.h"
 
 #define N_ROWS 8
-#define LG_WG_SIZE 9
+#define LG_WG_SIZE (7 + LG_WG_FACTOR)
 #define WG_SIZE (1 << LG_WG_SIZE)
 #define PARTITION_SIZE (WG_SIZE * N_ROWS)
 
diff --git a/piet-gpu/shader/transform_scan.comp b/piet-gpu/shader/transform_scan.comp
index 492bf04..c4d6745 100644
--- a/piet-gpu/shader/transform_scan.comp
+++ b/piet-gpu/shader/transform_scan.comp
@@ -3,9 +3,12 @@
 // A scan for a tree reduction prefix scan (either root or not, by ifdef).
 
 #version 450
+#extension GL_GOOGLE_include_directive : enable
+
+#include "setup.h"
 
 #define N_ROWS 8
-#define LG_WG_SIZE 9
+#define LG_WG_SIZE (7 + LG_WG_FACTOR)
 #define WG_SIZE (1 << LG_WG_SIZE)
 #define PARTITION_SIZE (WG_SIZE * N_ROWS)
 
diff --git a/piet-gpu/src/encoder.rs b/piet-gpu/src/encoder.rs
index fb32f26..87fff1c 100644
--- a/piet-gpu/src/encoder.rs
+++ b/piet-gpu/src/encoder.rs
@@ -19,7 +19,9 @@
 use bytemuck::{Pod, Zeroable};
 use piet_gpu_hal::BufWrite;
 
-use crate::stages::{self, Config, PathEncoder, Transform};
+use crate::stages::{
+    self, Config, PathEncoder, Transform, DRAW_PART_SIZE, PATHSEG_PART_SIZE, TRANSFORM_PART_SIZE,
+};
 
 pub struct Encoder {
     transform_stream: Vec<stages::Transform>,
@@ -52,12 +54,6 @@ const BBOX_SIZE: usize = 24;
 const DRAWMONOID_SIZE: usize = 8;
 const ANNOTATED_SIZE: usize = 40;
 
-// Maybe pull these from the relevant stages? In any case, they may depend
-// on runtime query of GPU (supported workgroup size).
-const TRANSFORM_PART_SIZE: usize = 4096;
-const PATHSEG_PART_SIZE: usize = 2048;
-const DRAWOBJ_PART_SIZE: usize = 4096;
-
 // These are bytemuck versions of elements currently defined in the
 // Element struct in piet-gpu-types; that's pretty much going away.
 
@@ -183,15 +179,15 @@ impl Encoder {
     pub fn stage_config(&self) -> (Config, usize) {
         // Layout of scene buffer
         let n_drawobj = self.n_drawobj();
-        let n_drawobj_padded = align_up(n_drawobj, DRAWOBJ_PART_SIZE);
+        let n_drawobj_padded = align_up(n_drawobj, DRAW_PART_SIZE as usize);
         let trans_offset = n_drawobj_padded * DRAWOBJ_SIZE;
         let n_trans = self.transform_stream.len();
-        let n_trans_padded = align_up(n_trans, TRANSFORM_PART_SIZE);
+        let n_trans_padded = align_up(n_trans, TRANSFORM_PART_SIZE as usize);
         let linewidth_offset = trans_offset + n_trans_padded * TRANSFORM_SIZE;
         let n_linewidth = self.linewidth_stream.len();
         let pathtag_offset = linewidth_offset + n_linewidth * LINEWIDTH_SIZE;
         let n_pathtag = self.tag_stream.len();
-        let n_pathtag_padded = align_up(n_pathtag, PATHSEG_PART_SIZE);
+        let n_pathtag_padded = align_up(n_pathtag, PATHSEG_PART_SIZE as usize);
         let pathseg_offset = pathtag_offset + n_pathtag_padded;
 
         // Layout of memory
@@ -230,14 +226,14 @@ impl Encoder {
     pub fn write_scene(&self, buf: &mut BufWrite) {
         buf.extend_slice(&self.drawobj_stream);
         let n_drawobj = self.drawobj_stream.len() / DRAWOBJ_SIZE;
-        buf.fill_zero(padding(n_drawobj, DRAWOBJ_PART_SIZE) * DRAWOBJ_SIZE);
+        buf.fill_zero(padding(n_drawobj, DRAW_PART_SIZE as usize) * DRAWOBJ_SIZE);
         buf.extend_slice(&self.transform_stream);
         let n_trans = self.transform_stream.len();
-        buf.fill_zero(padding(n_trans, TRANSFORM_PART_SIZE) * TRANSFORM_SIZE);
+        buf.fill_zero(padding(n_trans, TRANSFORM_PART_SIZE as usize) * TRANSFORM_SIZE);
         buf.extend_slice(&self.linewidth_stream);
         buf.extend_slice(&self.tag_stream);
         let n_pathtag = self.tag_stream.len();
-        buf.fill_zero(padding(n_pathtag, PATHSEG_PART_SIZE));
+        buf.fill_zero(padding(n_pathtag, PATHSEG_PART_SIZE as usize));
         buf.extend_slice(&self.pathseg_stream);
     }
 
diff --git a/piet-gpu/src/render_ctx.rs b/piet-gpu/src/render_ctx.rs
index 31fbf9e..96bbf03 100644
--- a/piet-gpu/src/render_ctx.rs
+++ b/piet-gpu/src/render_ctx.rs
@@ -11,7 +11,7 @@ use piet::{
 
 use piet_gpu_hal::BufWrite;
 use piet_gpu_types::encoder::{Encode, Encoder};
-use piet_gpu_types::scene::{Element, SetFillMode};
+use piet_gpu_types::scene::Element;
 
 use crate::gradient::{LinearGradient, RampCache};
 use crate::text::Font;
@@ -25,7 +25,6 @@ pub struct PietGpuRenderContext {
     // Will probably need direct accesss to hal Device to create images etc.
     inner_text: PietGpuText,
     stroke_width: f32,
-    fill_mode: FillMode,
     // We're tallying these cpu-side for expedience, but will probably
     // move this to some kind of readback from element processing.
     /// The count of elements that make it through to coarse rasterization.
@@ -69,14 +68,6 @@ struct ClipElement {
     bbox: Option<Rect>,
 }
 
-#[derive(Clone, Copy, PartialEq)]
-pub(crate) enum FillMode {
-    // Fill path according to the non-zero winding rule.
-    Nonzero = 0,
-    // Fill stroked path.
-    Stroke = 1,
-}
-
 const TOLERANCE: f64 = 0.25;
 
 impl PietGpuRenderContext {
@@ -91,7 +82,6 @@ impl PietGpuRenderContext {
             elements,
             inner_text,
             stroke_width,
-            fill_mode: FillMode::Nonzero,
             path_count: 0,
             pathseg_count: 0,
             trans_count: 0,
@@ -160,15 +150,6 @@ impl PietGpuRenderContext {
     pub fn get_ramp_data(&self) -> Vec<u32> {
         self.ramp_cache.get_ramp_data()
     }
-
-    pub(crate) fn set_fill_mode(&mut self, fill_mode: FillMode) {
-        if self.fill_mode != fill_mode {
-            self.elements.push(Element::SetFillMode(SetFillMode {
-                fill_mode: fill_mode as u32,
-            }));
-            self.fill_mode = fill_mode;
-        }
-    }
 }
 
 impl RenderContext for PietGpuRenderContext {
diff --git a/piet-gpu/src/stages.rs b/piet-gpu/src/stages.rs
index 1683cac..014cef4 100644
--- a/piet-gpu/src/stages.rs
+++ b/piet-gpu/src/stages.rs
@@ -22,10 +22,12 @@ mod transform;
 
 use bytemuck::{Pod, Zeroable};
 
-pub use draw::{DrawBinding, DrawCode, DrawMonoid, DrawStage};
-pub use path::{PathBinding, PathCode, PathEncoder, PathStage};
+pub use draw::{DrawBinding, DrawCode, DrawMonoid, DrawStage, DRAW_PART_SIZE};
+pub use path::{PathBinding, PathCode, PathEncoder, PathStage, PATHSEG_PART_SIZE};
 use piet_gpu_hal::{Buffer, CmdBuf, Session};
-pub use transform::{Transform, TransformBinding, TransformCode, TransformStage};
+pub use transform::{
+    Transform, TransformBinding, TransformCode, TransformStage, TRANSFORM_PART_SIZE,
+};
 
 /// The configuration block passed to piet-gpu shaders.
 ///
diff --git a/piet-gpu/src/stages/draw.rs b/piet-gpu/src/stages/draw.rs
index da773cf..5328a84 100644
--- a/piet-gpu/src/stages/draw.rs
+++ b/piet-gpu/src/stages/draw.rs
@@ -30,9 +30,9 @@ pub struct DrawMonoid {
     pub clip_ix: u32,
 }
 
-const DRAW_WG: u64 = 512;
+const DRAW_WG: u64 = 256;
 const DRAW_N_ROWS: u64 = 8;
-const DRAW_PART_SIZE: u64 = DRAW_WG * DRAW_N_ROWS;
+pub const DRAW_PART_SIZE: u64 = DRAW_WG * DRAW_N_ROWS;
 
 pub struct DrawCode {
     reduce_pipeline: Pipeline,
diff --git a/piet-gpu/src/stages/path.rs b/piet-gpu/src/stages/path.rs
index c9d2c60..6c524a2 100644
--- a/piet-gpu/src/stages/path.rs
+++ b/piet-gpu/src/stages/path.rs
@@ -39,22 +39,23 @@ pub struct PathBinding {
 }
 
 const REDUCE_WG: u32 = 128;
-const REDUCE_N_ROWS: u32 = 4;
+const REDUCE_N_ROWS: u32 = 2;
 const REDUCE_PART_SIZE: u32 = REDUCE_WG * REDUCE_N_ROWS;
 
-const ROOT_WG: u32 = 512;
+const ROOT_WG: u32 = 256;
 const ROOT_N_ROWS: u32 = 8;
 const ROOT_PART_SIZE: u32 = ROOT_WG * ROOT_N_ROWS;
 
-const SCAN_WG: u32 = 512;
+const SCAN_WG: u32 = 256;
 const SCAN_N_ROWS: u32 = 4;
 const SCAN_PART_SIZE: u32 = SCAN_WG * SCAN_N_ROWS;
 
-const CLEAR_WG: u32 = 512;
+pub const PATHSEG_PART_SIZE: u32 = SCAN_PART_SIZE;
+
+const CLEAR_WG: u32 = 256;
 
 impl PathCode {
     pub unsafe fn new(session: &Session) -> PathCode {
-        // TODO: add cross-compilation
         let reduce_code = include_shader!(session, "../../shader/gen/pathtag_reduce");
         let reduce_pipeline = session
             .create_compute_pipeline(
diff --git a/piet-gpu/src/stages/transform.rs b/piet-gpu/src/stages/transform.rs
index 4383c14..b21712f 100644
--- a/piet-gpu/src/stages/transform.rs
+++ b/piet-gpu/src/stages/transform.rs
@@ -33,9 +33,9 @@ pub struct Transform {
     pub translate: [f32; 2],
 }
 
-const TRANSFORM_WG: u64 = 512;
+const TRANSFORM_WG: u64 = 256;
 const TRANSFORM_N_ROWS: u64 = 8;
-const TRANSFORM_PART_SIZE: u64 = TRANSFORM_WG * TRANSFORM_N_ROWS;
+pub const TRANSFORM_PART_SIZE: u64 = TRANSFORM_WG * TRANSFORM_N_ROWS;
 
 pub struct TransformCode {
     reduce_pipeline: Pipeline,
diff --git a/piet-gpu/src/text.rs b/piet-gpu/src/text.rs
index a47c614..dec3ffa 100644
--- a/piet-gpu/src/text.rs
+++ b/piet-gpu/src/text.rs
@@ -11,7 +11,7 @@ use piet::{
 };
 
 use crate::encoder::GlyphEncoder;
-use crate::render_ctx::{self, FillMode};
+use crate::render_ctx;
 use crate::stages::Transform;
 use crate::PietGpuRenderContext;
 
@@ -172,7 +172,6 @@ impl PietGpuTextLayout {
         let mut inv_transform = None;
         // TODO: handle y offsets also
         let mut last_x = 0.0;
-        ctx.set_fill_mode(FillMode::Nonzero);
         for glyph in &self.glyphs {
             let transform = match &mut inv_transform {
                 None => {
diff --git a/tests/src/draw.rs b/tests/src/draw.rs
index 916f14c..d79a9d9 100644
--- a/tests/src/draw.rs
+++ b/tests/src/draw.rs
@@ -38,7 +38,8 @@ struct DrawTestData {
 
 pub unsafe fn draw_test(runner: &mut Runner, config: &Config) -> TestResult {
     let mut result = TestResult::new("draw");
-    let n_tag: u64 = config.size.choose(1 << 12, 1 << 20, 1 << 24);
+    // TODO: implement large scan and set large to 1 << 24
+    let n_tag: u64 = config.size.choose(1 << 12, 1 << 20, 1 << 22);
     let data = DrawTestData::new(n_tag);
     let stage_config = data.get_config();
 
diff --git a/tests/src/path.rs b/tests/src/path.rs
index 7c5388f..6f1f61a 100644
--- a/tests/src/path.rs
+++ b/tests/src/path.rs
@@ -62,7 +62,8 @@ struct Bbox {
 pub unsafe fn path_test(runner: &mut Runner, config: &Config) -> TestResult {
     let mut result = TestResult::new("path");
 
-    let n_path: u64 = config.size.choose(1 << 12, 1 << 16, 1 << 18);
+    // TODO: implement large scans and raise limit
+    let n_path: u64 = config.size.choose(1 << 12, 1 << 16, 209_000);
     let path_data = PathData::new(n_path as u32);
     let stage_config = path_data.get_config();
     let config_buf = runner
diff --git a/tests/src/transform.rs b/tests/src/transform.rs
index 1c15634..6edcc3f 100644
--- a/tests/src/transform.rs
+++ b/tests/src/transform.rs
@@ -30,7 +30,8 @@ struct AffineTestData {
 
 pub unsafe fn transform_test(runner: &mut Runner, config: &Config) -> TestResult {
     let mut result = TestResult::new("transform");
-    let n_elements: u64 = config.size.choose(1 << 12, 1 << 18, 1 << 24);
+    // TODO: implement large scan and set large to 1 << 24
+    let n_elements: u64 = config.size.choose(1 << 12, 1 << 18, 1 << 22);
     // Validate with real transform data.
     let data = AffineTestData::new(n_elements as usize);
     let data_buf = runner