mirror of
https://github.com/italicsjenga/vello.git
synced 2024-10-16 22:51:31 +11:00
1dea6c0ef0
Fixed several other shader validation errors caught when running vello_shaders natively on Metal. These were primarily caused by reading an invalid drawtag while accessing the scene buffer. Scene buffer access in the offending pipelines now initialize the draw tag to DRAWTAG_NOP if an invocation ID would land beyond the valid index range of encoded draw objects.
43 lines
1.1 KiB
WebGPU Shading Language
43 lines
1.1 KiB
WebGPU Shading Language
// SPDX-License-Identifier: Apache-2.0 OR MIT OR Unlicense
|
|
|
|
#import config
|
|
#import drawtag
|
|
|
|
@group(0) @binding(0)
|
|
var<uniform> config: Config;
|
|
|
|
@group(0) @binding(1)
|
|
var<storage> scene: array<u32>;
|
|
|
|
@group(0) @binding(2)
|
|
var<storage, read_write> reduced: array<DrawMonoid>;
|
|
|
|
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 = 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) {
|
|
workgroupBarrier();
|
|
if local_id.x + (1u << i) < WG_SIZE {
|
|
let other = sh_scratch[local_id.x + (1u << i)];
|
|
agg = combine_draw_monoid(agg, other);
|
|
}
|
|
workgroupBarrier();
|
|
sh_scratch[local_id.x] = agg;
|
|
}
|
|
if local_id.x == 0u {
|
|
reduced[ix >> firstTrailingBit(WG_SIZE)] = agg;
|
|
}
|
|
}
|