diff --git a/piet-scene/src/glyph.rs b/piet-scene/src/glyph.rs
index d18a89e..8113af4 100644
--- a/piet-scene/src/glyph.rs
+++ b/piet-scene/src/glyph.rs
@@ -124,12 +124,14 @@ impl<'a> GlyphProvider<'a> {
                     if let Some(xform) = xform_stack.last() {
                         builder.push_layer(
                             Mix::Clip,
+                            1.0,
                             Affine::IDENTITY,
                             &convert_transformed_path(path.elements(), xform),
                         );
                     } else {
                         builder.push_layer(
                             Mix::Clip,
+                            1.0,
                             Affine::IDENTITY,
                             &convert_path(path.elements()),
                         );
@@ -144,7 +146,7 @@ impl<'a> GlyphProvider<'a> {
                         max = *xform * max;
                     }
                     let rect = Rect::from_points(min, max);
-                    builder.push_layer(Mix::Normal, Affine::IDENTITY, &rect);
+                    builder.push_layer(Mix::Normal, 1.0, Affine::IDENTITY, &rect);
                 }
                 Command::PopLayer => builder.pop_layer(),
                 Command::BeginBlend(bounds, mode) => {
@@ -155,7 +157,7 @@ impl<'a> GlyphProvider<'a> {
                         max = *xform * max;
                     }
                     let rect = Rect::from_points(min, max);
-                    builder.push_layer(convert_blend(*mode), Affine::IDENTITY, &rect);
+                    builder.push_layer(convert_blend(*mode), 1.0, Affine::IDENTITY, &rect);
                 }
                 Command::EndBlend => builder.pop_layer(),
                 Command::SimpleFill(path_index, brush, brush_xform) => {
diff --git a/piet-scene/src/scene/builder.rs b/piet-scene/src/scene/builder.rs
index 17a236a..b202cea 100644
--- a/piet-scene/src/scene/builder.rs
+++ b/piet-scene/src/scene/builder.rs
@@ -24,7 +24,7 @@ use smallvec::SmallVec;
 /// Builder for constructing a scene or scene fragment.
 pub struct SceneBuilder<'a> {
     scene: &'a mut SceneData,
-    layers: SmallVec<[BlendMode; 8]>,
+    layer_depth: u32,
 }
 
 impl<'a> SceneBuilder<'a> {
@@ -45,7 +45,7 @@ impl<'a> SceneBuilder<'a> {
         scene.reset(is_fragment);
         Self {
             scene,
-            layers: Default::default(),
+            layer_depth: 0,
         }
     }
 
@@ -54,6 +54,7 @@ impl<'a> SceneBuilder<'a> {
     pub fn push_layer(
         &mut self,
         blend: impl Into<BlendMode>,
+        alpha: f32,
         transform: Affine,
         shape: &impl Shape,
     ) {
@@ -65,14 +66,15 @@ impl<'a> SceneBuilder<'a> {
             // all drawing until the layer is popped.
             self.encode_path(&Rect::new(0.0, 0.0, 0.0, 0.0), true);
         }
-        self.begin_clip(blend);
-        self.layers.push(blend);
+        self.begin_clip(blend, alpha.clamp(0.0, 1.0));
+        self.layer_depth += 1;
     }
 
     /// Pops the current layer.
     pub fn pop_layer(&mut self) {
-        if let Some(blend) = self.layers.pop() {
-            self.end_clip(blend);
+        if self.layer_depth > 0 {
+            self.end_clip();
+            self.layer_depth -= 1;
         }
     }
 
@@ -127,8 +129,8 @@ impl<'a> SceneBuilder<'a> {
 
     /// Completes construction and finalizes the underlying scene.
     pub fn finish(mut self) {
-        while !self.layers.is_empty() {
-            self.pop_layer();
+        for _ in 0..self.layer_depth {
+            self.end_clip();
         }
     }
 }
@@ -250,10 +252,11 @@ impl<'a> SceneBuilder<'a> {
     }
 
     /// Start a clip.
-    fn begin_clip(&mut self, blend: BlendMode) {
+    fn begin_clip(&mut self, blend: BlendMode, alpha: f32) {
         self.scene.drawtag_stream.push(DRAWTAG_BEGINCLIP);
         let element = Clip {
             blend: encode_blend_mode(blend),
+            alpha,
         };
         self.scene
             .drawdata_stream
@@ -261,14 +264,8 @@ impl<'a> SceneBuilder<'a> {
         self.scene.n_clip += 1;
     }
 
-    fn end_clip(&mut self, blend: BlendMode) {
+    fn end_clip(&mut self) {
         self.scene.drawtag_stream.push(DRAWTAG_ENDCLIP);
-        let element = Clip {
-            blend: encode_blend_mode(blend),
-        };
-        self.scene
-            .drawdata_stream
-            .extend(bytemuck::bytes_of(&element));
         // This is a dummy path, and will go away with the new clip impl.
         self.scene.tag_stream.push(0x10);
         self.scene.n_path += 1;
@@ -284,8 +281,8 @@ fn encode_blend_mode(mode: BlendMode) -> u32 {
 const DRAWTAG_FILLCOLOR: u32 = 0x44;
 const DRAWTAG_FILLLINGRADIENT: u32 = 0x114;
 const DRAWTAG_FILLRADGRADIENT: u32 = 0x2dc;
-const DRAWTAG_BEGINCLIP: u32 = 0x05;
-const DRAWTAG_ENDCLIP: u32 = 0x25;
+const DRAWTAG_BEGINCLIP: u32 = 0x89;
+const DRAWTAG_ENDCLIP: u32 = 0xa1;
 
 #[repr(C)]
 #[derive(Clone, Copy, Debug, Default, Zeroable, Pod)]
@@ -324,6 +321,7 @@ pub struct FillImage {
 #[derive(Clone, Copy, Debug, Default, Zeroable, Pod)]
 pub struct Clip {
     blend: u32,
+    alpha: f32,
 }
 
 struct PathBuilder<'a> {
diff --git a/piet-wgsl/examples/winit/src/test_scene.rs b/piet-wgsl/examples/winit/src/test_scene.rs
index 4c90829..047f286 100644
--- a/piet-wgsl/examples/winit/src/test_scene.rs
+++ b/piet-wgsl/examples/winit/src/test_scene.rs
@@ -146,7 +146,7 @@ fn render_clip_test(sb: &mut SceneBuilder) {
             PathEl::LineTo((X0, Y1).into()),
             PathEl::ClosePath,
         ];
-        sb.push_layer(Mix::Clip, Affine::IDENTITY, &path);
+        sb.push_layer(Mix::Clip, 1.0, Affine::IDENTITY, &path);
     }
     let rect = Rect::new(X0, Y0, X1, Y1);
     sb.fill(
@@ -178,7 +178,12 @@ fn render_alpha_test(sb: &mut SceneBuilder) {
         None,
         &make_diamond(1024.0, 125.0),
     );
-    sb.push_layer(Mix::Clip, Affine::IDENTITY, &make_diamond(1024.0, 150.0));
+    sb.push_layer(
+        Mix::Clip,
+        1.0,
+        Affine::IDENTITY,
+        &make_diamond(1024.0, 150.0),
+    );
     sb.fill(
         Fill::NonZero,
         Affine::IDENTITY,
@@ -240,10 +245,10 @@ fn render_blend_square(sb: &mut SceneBuilder, blend: BlendMode, transform: Affin
         Color::rgb8(0, 255, 0),
         Color::rgb8(0, 0, 255),
     ];
-    sb.push_layer(Mix::Normal, transform, &rect);
+    sb.push_layer(Mix::Normal, 1.0, transform, &rect);
     for (i, c) in COLORS.iter().enumerate() {
         let linear = LinearGradient::new((0.0, 0.0), (0.0, 200.0)).stops([Color::WHITE, *c]);
-        sb.push_layer(blend, transform, &rect);
+        sb.push_layer(blend, 1.0, transform, &rect);
         // squash the ellipse
         let a = transform
             * Affine::translate((100., 100.))
@@ -273,12 +278,13 @@ fn blend_square(blend: BlendMode) -> SceneFragment {
 
 #[allow(unused)]
 pub fn render_anim_frame(sb: &mut SceneBuilder, text: &mut SimpleText, i: usize) {
+    let rect = Rect::from_origin_size(Point::new(0.0, 0.0), (1000.0, 1000.0));
     sb.fill(
         Fill::NonZero,
         Affine::IDENTITY,
         &Brush::Solid(Color::rgb8(128, 128, 128)),
         None,
-        &Rect::from_origin_size(Point::new(0.0, 0.0), (1000.0, 1000.0)),
+        &rect,
     );
     let text_size = 60.0 + 40.0 * (0.01 * i as f32).sin();
     let s = "\u{1f600}hello piet-gpu text!";
@@ -308,8 +314,32 @@ pub fn render_anim_frame(sb: &mut SceneBuilder, text: &mut SimpleText, i: usize)
         Affine::IDENTITY,
         &Brush::Solid(Color::rgb8(128, 0, 0)),
         None,
-        &&[PathEl::MoveTo(center), PathEl::LineTo(p1)][..],
+        &[PathEl::MoveTo(center), PathEl::LineTo(p1)],
     );
+    sb.fill(
+        Fill::NonZero,
+        Affine::translate((150.0, 150.0)) * Affine::scale(0.2),
+        Color::RED,
+        None,
+        &rect,
+    );
+    let alpha = (i as f64 * 0.03).sin() as f32 * 0.5 + 0.5;
+    sb.push_layer(Mix::Normal, alpha, Affine::IDENTITY, &rect);
+    sb.fill(
+        Fill::NonZero,
+        Affine::translate((100.0, 100.0)) * Affine::scale(0.2),
+        Color::BLUE,
+        None,
+        &rect,
+    );
+    sb.fill(
+        Fill::NonZero,
+        Affine::translate((200.0, 200.0)) * Affine::scale(0.2),
+        Color::GREEN,
+        None,
+        &rect,
+    );
+    sb.pop_layer();
 }
 
 #[allow(unused)]
diff --git a/piet-wgsl/shader/clip_leaf.wgsl b/piet-wgsl/shader/clip_leaf.wgsl
index db4cb6d..219ba76 100644
--- a/piet-wgsl/shader/clip_leaf.wgsl
+++ b/piet-wgsl/shader/clip_leaf.wgsl
@@ -9,7 +9,7 @@
 var<storage> config: Config;
 
 @group(0) @binding(1)
-var<storage> clip_inp: array<i32>;
+var<storage> clip_inp: array<ClipInp>;
 
 @group(0) @binding(2)
 var<storage> path_bboxes: array<PathBbox>;
@@ -26,6 +26,9 @@ var<storage, read_write> draw_monoids: array<DrawMonoid>;
 @group(0) @binding(6)
 var<storage, read_write> clip_bboxes: array<vec4<f32>>;
 
+@group(0) @binding(7)
+var<storage, read_write> info: array<u32>;
+
 let WG_SIZE = 256u;
 var<workgroup> sh_bic: array<Bic, 510 >;
 var<workgroup> sh_stack: array<u32, WG_SIZE>;
@@ -68,7 +71,7 @@ fn search_link(bic: ptr<function, Bic>, ix: u32) -> i32 {
 
 fn load_clip_inp(ix: u32) -> i32 {
     if ix < config.n_clip {
-        return clip_inp[ix];
+        return clip_inp[ix].path_ix;
     } else {
         return -2147483648;
         // literal too large?
@@ -185,8 +188,16 @@ fn main(
 
     if !is_push && global_id.x < config.n_clip {
         // Fix up drawmonoid so path_ix of EndClip matches BeginClip
-        let path_ix = clip_inp[parent];
-        draw_monoids[~inp].path_ix = u32(path_ix);
+        let parent_clip = clip_inp[parent];
+        let path_ix = parent_clip.path_ix;
+        let parent_ix = parent_clip.ix;
+        let ix = ~inp;
+        draw_monoids[ix].path_ix = u32(path_ix);
+        // Copy blend mode and alpha from parent
+        let di = draw_monoids[ix].info_offset;
+        let parent_di = draw_monoids[parent_ix].info_offset;
+        info[di] = info[parent_di];
+        info[di + 1u] = info[parent_di + 1u];
 
         if grandparent >= 0 {
             bbox = sh_bbox[grandparent];
diff --git a/piet-wgsl/shader/clip_reduce.wgsl b/piet-wgsl/shader/clip_reduce.wgsl
index b314fe8..2b5b60c 100644
--- a/piet-wgsl/shader/clip_reduce.wgsl
+++ b/piet-wgsl/shader/clip_reduce.wgsl
@@ -8,7 +8,7 @@
 var<storage> config: Config;
 
 @group(0) @binding(1)
-var<storage> clip_inp: array<i32>;
+var<storage> clip_inp: array<ClipInp>;
 
 @group(0) @binding(2)
 var<storage> path_bboxes: array<PathBbox>;
@@ -30,7 +30,7 @@ fn main(
     @builtin(local_invocation_id) local_id: vec3<u32>,
     @builtin(workgroup_id) wg_id: vec3<u32>,
 ) {
-    let inp = clip_inp[global_id.x];
+    let inp = clip_inp[global_id.x].path_ix;
     let is_push = inp >= 0;
     var bic = Bic(1u - u32(is_push), u32(is_push));
     // reverse scan of bicyclic semigroup
diff --git a/piet-wgsl/shader/coarse.wgsl b/piet-wgsl/shader/coarse.wgsl
index 2f78c5c..2bd14d8 100644
--- a/piet-wgsl/shader/coarse.wgsl
+++ b/piet-wgsl/shader/coarse.wgsl
@@ -127,11 +127,12 @@ fn write_begin_clip() {
     cmd_offset += 1u;
 }
 
-fn write_end_clip(blend: u32) {
-    alloc_cmd(2u);
+fn write_end_clip(end_clip: CmdEndClip) {
+    alloc_cmd(3u);
     ptcl[cmd_offset] = CMD_END_CLIP;
-    ptcl[cmd_offset + 1u] = blend;
-    cmd_offset += 2u;
+    ptcl[cmd_offset + 1u] = end_clip.blend;
+    ptcl[cmd_offset + 2u] = bitcast<u32>(end_clip.alpha);
+    cmd_offset += 3u;            
 }
 
 @compute @workgroup_size(256)
@@ -270,7 +271,6 @@ fn main(
             }
             drawobj_ix = sh_drawobj_ix[el_ix];
             tag = scene[config.drawtag_base + drawobj_ix];
-            // TODO: clip logic
             let seq_ix = ix - select(0u, sh_tile_count[el_ix - 1u], el_ix > 0u);
             let width = sh_tile_width[el_ix];
             let x = sh_tile_x0[el_ix] + seq_ix % width;
@@ -281,9 +281,8 @@ fn main(
             var is_blend = false;
             if is_clip {
                 let BLEND_CLIP = (128u << 8u) | 3u;
-                let scene_offset = draw_monoids[drawobj_ix].scene_offset;
-                let dd = config.drawdata_base + scene_offset;
-                let blend = scene[dd];
+                let di = draw_monoids[drawobj_ix].info_offset;
+                let blend = info[di];
                 is_blend = blend != BLEND_CLIP;
             }
             let include_tile = tile.segments != 0u || (tile.backdrop == 0) == is_clip || is_blend;
@@ -348,7 +347,7 @@ fn main(
                         write_grad(CMD_RAD_GRAD, index, info_offset);
                     }
                     // DRAWTAG_BEGIN_CLIP
-                    case 0x05u: {
+                    case 0x89u: {
                         if tile.segments == 0u && tile.backdrop == 0 {
                             clip_zero_depth = clip_depth + 1u;
                         } else {
@@ -359,10 +358,12 @@ fn main(
                         clip_depth += 1u;
                     }
                     // DRAWTAG_END_CLIP
-                    case 0x25u: {
+                    case 0xa1u: {
                         clip_depth -= 1u;
                         write_path(tile, -1.0);
-                        write_end_clip(scene[dd]);
+                        let blend = info[di];
+                        let alpha = bitcast<f32>(info[di + 1u]);
+                        write_end_clip(CmdEndClip(blend, alpha));
                         render_blend_depth -= 1u;
                     }
                     default: {}
@@ -371,11 +372,11 @@ fn main(
                 // In "clip zero" state, suppress all drawing
                 switch drawtag {
                     // DRAWTAG_BEGIN_CLIP
-                    case 0x05u: {
+                    case 0x89u: {
                         clip_depth += 1u;
                     }
                     // DRAWTAG_END_CLIP
-                    case 0x25u: {
+                    case 0xa1u: {
                         if clip_depth == clip_zero_depth {
                             clip_zero_depth = 0u;
                         }
diff --git a/piet-wgsl/shader/draw_leaf.wgsl b/piet-wgsl/shader/draw_leaf.wgsl
index cd3f843..05bd80f 100644
--- a/piet-wgsl/shader/draw_leaf.wgsl
+++ b/piet-wgsl/shader/draw_leaf.wgsl
@@ -26,7 +26,7 @@ var<storage, read_write> draw_monoid: array<DrawMonoid>;
 var<storage, read_write> info: array<u32>;
 
 @group(0) @binding(6)
-var<storage, read_write> clip_inp: array<i32>;
+var<storage, read_write> clip_inp: array<ClipInp>;
 
 let WG_SIZE = 256u;
 
@@ -170,6 +170,15 @@ fn main(
                 info[di + 9u] = bitcast<u32>(ra);
                 info[di + 10u] = bitcast<u32>(roff);
             }
+            // DRAWTAG_BEGIN_CLIP
+            case 0x89u: {
+                // Store blend mode and alpha in info for two reasons: 1) we don't need
+                // to bind scene in clip_leaf which keeps us at 8 buffer bindings and 2)
+                // the logic in coarse to check clip state for tile inclusion is the
+                // same for BeginClip/EndClip.
+                info[di] = scene[dd];
+                info[di + 1u] = scene[dd + 1u];
+            }
             default: {}
         }
     }
@@ -178,6 +187,6 @@ fn main(
         if tag_word == DRAWTAG_BEGIN_CLIP {
             path_ix = m.path_ix;
         }
-        clip_inp[m.clip_ix] = i32(path_ix);
+        clip_inp[m.clip_ix] = ClipInp(ix, i32(path_ix));
     }
 }
diff --git a/piet-wgsl/shader/fine.wgsl b/piet-wgsl/shader/fine.wgsl
index 85d9950..6c23035 100644
--- a/piet-wgsl/shader/fine.wgsl
+++ b/piet-wgsl/shader/fine.wgsl
@@ -82,6 +82,12 @@ fn read_rad_grad(cmd_ix: u32) -> CmdRadGrad {
     return CmdRadGrad(index, matrx, xlat, c1, ra, roff);
 }
 
+fn read_end_clip(cmd_ix: u32) -> CmdEndClip {
+    let blend = ptcl[cmd_ix + 1u];
+    let alpha = bitcast<f32>(ptcl[cmd_ix + 2u]);
+    return CmdEndClip(blend, alpha);
+}
+
 #else
 
 @group(0) @binding(3)
@@ -260,7 +266,7 @@ fn main(
             }
             // CMD_END_CLIP
             case 10u: {
-                let blend = ptcl[cmd_ix + 1u];
+                let end_clip = read_end_clip(cmd_ix);
                 clip_depth -= 1u;
                 for (var i = 0u; i < PIXELS_PER_THREAD; i += 1u) {
                     var bg_rgba: u32;
@@ -270,10 +276,10 @@ fn main(
                         // load from memory
                     }
                     let bg = unpack4x8unorm(bg_rgba);
-                    let fg = rgba[i] * area[i];
-                    rgba[i] = blend_mix_compose(bg, fg, blend);
+                    let fg = rgba[i] * area[i] * end_clip.alpha;
+                    rgba[i] = blend_mix_compose(bg, fg, end_clip.blend);
                 }
-                cmd_ix += 2u;
+                cmd_ix += 3u;
             }
             // CMD_JUMP
             case 11u: {
diff --git a/piet-wgsl/shader/shared/clip.wgsl b/piet-wgsl/shader/shared/clip.wgsl
index e2b9fd6..a845ce9 100644
--- a/piet-wgsl/shader/shared/clip.wgsl
+++ b/piet-wgsl/shader/shared/clip.wgsl
@@ -10,6 +10,11 @@ fn bic_combine(x: Bic, y: Bic) -> Bic {
     return Bic(x.a + y.a - m, x.b + y.b - m);
 }
 
+struct ClipInp {
+    ix: u32,
+    path_ix: i32,
+}
+
 struct ClipEl {
     parent_ix: u32,
     bbox: vec4<f32>,
diff --git a/piet-wgsl/shader/shared/drawtag.wgsl b/piet-wgsl/shader/shared/drawtag.wgsl
index 93554a7..3b68c55 100644
--- a/piet-wgsl/shader/shared/drawtag.wgsl
+++ b/piet-wgsl/shader/shared/drawtag.wgsl
@@ -20,8 +20,8 @@ let DRAWTAG_FILL_COLOR = 0x44u;
 let DRAWTAG_FILL_LIN_GRADIENT = 0x114u;
 let DRAWTAG_FILL_RAD_GRADIENT = 0x2dcu;
 let DRAWTAG_FILL_IMAGE = 0x48u;
-let DRAWTAG_BEGIN_CLIP = 0x05u;
-let DRAWTAG_END_CLIP = 0x25u;
+let DRAWTAG_BEGIN_CLIP = 0x89u;
+let DRAWTAG_END_CLIP = 0xa1u;
 
 fn draw_monoid_identity() -> DrawMonoid {
     return DrawMonoid();
diff --git a/piet-wgsl/shader/shared/ptcl.wgsl b/piet-wgsl/shader/shared/ptcl.wgsl
index a080f2f..3527bd0 100644
--- a/piet-wgsl/shader/shared/ptcl.wgsl
+++ b/piet-wgsl/shader/shared/ptcl.wgsl
@@ -56,3 +56,8 @@ struct CmdRadGrad {
     ra: f32,
     roff: f32,
 }
+
+struct CmdEndClip {
+    blend: u32,
+    alpha: f32,
+}
diff --git a/piet-wgsl/src/render.rs b/piet-wgsl/src/render.rs
index 2d50881..ff8664d 100644
--- a/piet-wgsl/src/render.rs
+++ b/piet-wgsl/src/render.rs
@@ -16,7 +16,7 @@ const DRAWMONOID_SIZE: u64 = 16;
 const MAX_DRAWINFO_SIZE: u64 = 44;
 const CLIP_BIC_SIZE: u64 = 8;
 const CLIP_EL_SIZE: u64 = 32;
-const CLIP_INP_SIZE: u64 = 4;
+const CLIP_INP_SIZE: u64 = 8;
 const CLIP_BBOX_SIZE: u64 = 16;
 const PATH_SIZE: u64 = 32;
 const DRAW_BBOX_SIZE: u64 = 16;
@@ -324,6 +324,7 @@ pub fn render_full(
                 clip_el_buf,
                 draw_monoid_buf,
                 clip_bbox_buf,
+                info_buf,
             ],
         );
     }
diff --git a/piet-wgsl/src/shaders.rs b/piet-wgsl/src/shaders.rs
index a8a4176..186340b 100644
--- a/piet-wgsl/src/shaders.rs
+++ b/piet-wgsl/src/shaders.rs
@@ -212,6 +212,7 @@ pub fn full_shaders(device: &Device, engine: &mut Engine) -> Result<FullShaders,
             BindType::BufReadOnly,
             BindType::Buffer,
             BindType::Buffer,
+            BindType::Buffer,
         ],
     )?;
     let binning = engine.add_shader(