Merge pull request #329 from armansito/draw-monoid-overrun

[draw_leaf] Guard against invalid reads of drawtags and bad writes to draw monoids.

.vscode/settings.json

@@ -12,6 +12,7 @@
     "segment": "${workspaceFolder}/shader/shared/segment.wgsl",
     "tile": "${workspaceFolder}/shader/shared/tile.wgsl",
     "transform": "${workspaceFolder}/shader/shared/transform.wgsl"
+    "util": "${workspaceFolder}/shader/shared/util.wgsl",
   },
   "wgsl-analyzer.diagnostics.nagaVersion": "main",
   "wgsl-analyzer.preprocessor.shaderDefs": [

shader/binning.wgsl

@@ -74,7 +74,10 @@ fn main(
         let draw_monoid = draw_monoids[element_ix];
         var clip_bbox = vec4(-1e9, -1e9, 1e9, 1e9);
         if draw_monoid.clip_ix > 0u {
-            clip_bbox = clip_bbox_buf[draw_monoid.clip_ix - 1u];
+            // TODO: `clip_ix` should always be valid as long as the monoids are correct. Leaving
+            // the bounds check in here for correctness but we should assert this condition instead
+            // once there is a debug-assertion mechanism.
+            clip_bbox = clip_bbox_buf[min(draw_monoid.clip_ix - 1u, config.n_clip - 1u)];
         }
         // For clip elements, clip_box is the bbox of the clip path,
         // intersected with enclosing clips.

shader/draw_leaf.wgsl

@@ -29,6 +29,8 @@ var<storage, read_write> info: array<u32>;
 @group(0) @binding(6)
 var<storage, read_write> clip_inp: array<ClipInp>;
 
+#import util
+
 let WG_SIZE = 256u;
 
 fn read_transform(transform_base: u32, ix: u32) -> Transform {
@@ -73,7 +75,7 @@ fn main(
     workgroupBarrier();
     var m = sh_scratch[0];
     workgroupBarrier();
-    let tag_word = scene[config.drawtag_base + ix];
+    let tag_word = read_draw_tag_from_scene(ix);
     agg = map_draw_tag(tag_word);
     sh_scratch[local_id.x] = agg;
     for (var i = 0u; i < firstTrailingBit(WG_SIZE); i += 1u) {
@@ -90,7 +92,9 @@ fn main(
         m = combine_draw_monoid(m, sh_scratch[local_id.x - 1u]);
     }
     // m now contains exclusive prefix sum of draw monoid
-    draw_monoid[ix] = m;
+    if ix < config.n_drawobj {
+        draw_monoid[ix] = m;
+    }
     let dd = config.drawdata_base + m.scene_offset;
     let di = m.info_offset;
     if tag_word == DRAWTAG_FILL_COLOR || tag_word == DRAWTAG_FILL_LIN_GRADIENT ||

shader/draw_reduce.wgsl

@@ -16,13 +16,15 @@ let WG_SIZE = 256u;
 
 var<workgroup> sh_scratch: array<DrawMonoid, WG_SIZE>;
 
+#import util
+
 @compute @workgroup_size(256)
 fn main(
     @builtin(global_invocation_id) global_id: vec3<u32>,
     @builtin(local_invocation_id) local_id: vec3<u32>,
 ) {
     let ix = global_id.x;
-    let tag_word = scene[config.drawtag_base + ix];
+    let tag_word = read_draw_tag_from_scene(ix);
     var agg = map_draw_tag(tag_word);
     sh_scratch[local_id.x] = agg;
     for (var i = 0u; i < firstTrailingBit(WG_SIZE); i += 1u) {

shader/shared/util.wgsl

@@ -0,0 +1,23 @@
+// SPDX-License-Identifier: Apache-2.0 OR MIT OR Unlicense
+
+// This file defines utility functions that interact with host-shareable buffer objects. It should
+// be imported once following the resource binding declarations in the shader module that access
+// them.
+
+// Reads a draw tag from the scene buffer, defaulting to DRAWTAG_NOP if the given `ix` is beyond the
+// range of valid draw objects (e.g this can happen if `ix` is derived from an invocation ID in a
+// workgroup that partially spans valid range).
+//
+// This function depends on the following global declarations:
+//    * `scene`: array<u32>
+//    * `config`: Config (see config.wgsl)
+fn read_draw_tag_from_scene(ix: u32) -> u32 {
+    let tag_ix = config.drawtag_base + ix;
+    var tag_word: u32;
+    if tag_ix < config.drawtag_base + config.n_drawobj {
+        tag_word = scene[tag_ix];
+    } else {
+        tag_word = DRAWTAG_NOP;
+    }
+    return tag_word;
+}

src/shaders.rs

@@ -319,4 +319,5 @@ const SHARED_SHADERS: &[(&str, &str)] = &[
     shared_shader!("segment"),
     shared_shader!("tile"),
     shared_shader!("transform"),
+    shared_shader!("util"),
 ];