diff --git a/shader/merge.wgsl b/shader/merge.wgsl
index 28b8350..d20c5c5 100644
--- a/shader/merge.wgsl
+++ b/shader/merge.wgsl
@@ -8,8 +8,14 @@
     path_ix: u32,
     x: u32,
     y: u32,
-    delta: i32,
-    // TODO: slope etc
+    p0: u32, // packed
+    p1: u32, // packed
+}
+
+fn unpack_point(p: u32) -> vec2f {
+    let x = f32(p & 0xffffu) * (1.0 / 8192.0);
+    let y = f32(p >> 16u) * (1.0 / 8192.0);
+    return vec2(x, y);
 }
 
 @group(0) @binding(0)
@@ -34,10 +40,15 @@
     return c;
 }
 
-fn mm_histogram(t: Minitile) -> u32 {
-    // TODO: get these from tile
-    let xmin = 0u;
-    let xmax = 4u;
+fn mt_delta(t: Minitile) -> i32 {
+    return i32((t.p1 >> 16u) == 0u) - i32((t.p0 >> 16u) == 0u);
+}
+
+fn mt_histogram(t: Minitile) -> u32 {
+    let x0 = f32(t.p0 & 0xffffu) * (1.0 / 8192.0);
+    let x1 = f32(t.p1 & 0xffffu) * (1.0 / 8192.0);
+    let xmin = u32(floor(min(x0, x1)));
+    let xmax = u32(ceil(max(x0, x1)));
     let rshift = (4u - (xmax - xmin)) * 8u;
     let lshift = xmin * 8u;
     return (0x01010101u >> rshift) << lshift;
@@ -54,10 +65,11 @@
 var<workgroup> sh_histo: array<u32, WG_SIZE>;
 var<workgroup> sh_seg_end: array<u32, WG_SIZE>;
 var<workgroup> sh_inclusive_cols: array<u32, WG_SIZE>;
+var<workgroup> sh_area: array<atomic<i32>, WG_SIZE>;
+var<workgroup> sh_carryover: array<i32, 4>;
 
 @compute @workgroup_size(256)
 fn main(
-    @builtin(global_invocation_id) global_id: vec3<u32>,
     @builtin(local_invocation_id) local_id: vec3<u32>,
     @builtin(workgroup_id) wg_id: vec3<u32>,
 ) {
@@ -65,19 +77,20 @@
     var first = false;
     var first_x = false;
     // predicate? or pad?
-    let tile = input[global_id.x];
-    if global_id.x != 0u {
-        let prev = input[global_id.x - 1u];
+    let global_ix = wg_id.x * WG_SIZE + local_id.x;
+    let tile = input[global_ix];
+    if global_ix != 0u {
+        let prev = input[global_ix - 1u];
         first = tile.path_ix != prev.path_ix || tile.y != prev.y;
         first_x = first || tile.x != prev.x;
     }
-    let winding = tile.delta;
+    let winding = mt_delta(tile);
     let n_strips = u32(first);
     let start = select(0u, local_id.x, first);
     let start_x = select(0u, local_id.x, first_x);
     var agg = MergeMonoid(winding, n_strips, start, start_x);
 
-    let local_histo = mm_histogram(tile);
+    let local_histo = mt_histogram(tile);
     var histo = local_histo;
     sh_mm[local_id.x] = agg;
     sh_histo[local_id.x] = histo;
@@ -159,5 +172,34 @@
             }
         }
         // at this point, lo should index our tile
+        // TODO: predicate on ix < total_cols?
+        let render_tile = input[wg_id.x * WG_SIZE + lo];
+        var alphas = 0u;
+        for (var y = 0u; y < 4u; y++) {
+            if tile_within_col == 0u {
+                atomicStore(&sh_area[local_id.x], 0);
+            }
+            workgroupBarrier();
+            var area_init = 0;
+            if local_id.x == 0u && block_ix != 0u {
+                area_init = sh_carryover[y];
+            }
+            let area = area_init; // TODO: compute from tile
+            atomicAdd(&sh_area[local_id.x - tile_within_col], area);
+            workgroupBarrier();
+            if tile_within_col == 0u {
+                let summed_area = atomicLoad(&sh_area[local_id.x]);
+                if seg_end == WG_SIZE - 1u {
+                    // TODO: only if last column
+                    sh_carryover[y] = summed_area;
+                }
+                let winding_area = sh_mm[seg_end].winding * 256;
+                let alpha_u8 = u32(min(abs(summed_area + winding_area), 255));
+                alphas = (alphas >> 8u) + (alpha_u8 << 24u);
+            }
+        }
+        if tile_within_col == 0u {
+            // TODO: store alphas
+        }
     }
 }