1override supports_indirect_first_instance: bool;
2override write_d3d12_special_constants: bool;
3
4struct MetadataEntry {
5    // bits 0..30 are an offset into `src`
6    // bit 31 signifies that we are validating an indexed draw
7    src_offset: u32,
8    // bits 0..30 are an offset into `dst`
9    // bit 30 is the most significant bit of `vertex_or_index_limit`
10    // bit 31 is the most significant bit of `instance_limit`
11    dst_offset: u32,
12    vertex_or_index_limit: u32,
13    instance_limit: u32,
14}
15
16struct MetadataRange {
17    start: u32,
18    count: u32,
19}
20var<push_constant> metadata_range: MetadataRange;
21
22@group(0) @binding(0)
23var<storage, read> metadata: array<MetadataEntry>;
24@group(1) @binding(0)
25var<storage, read> src: array<u32>;
26@group(2) @binding(0)
27var<storage, read_write> dst: array<u32>;
28
29fn is_bit_set(data: u32, index: u32) -> bool {
30    return ((data >> index) & 1u) == 1u;
31}
32
33@compute @workgroup_size(64)
34fn main(@builtin(global_invocation_id) global_invocation_id: vec3u) {
35    if global_invocation_id.x >= metadata_range.count { return; }
36
37    let metadata = metadata[metadata_range.start + global_invocation_id.x];
38    var failed = false;
39
40    let is_indexed = is_bit_set(metadata.src_offset, 31);
41    let src_base_offset = ((metadata.src_offset << 2) >> 2);
42    let dst_base_offset = ((metadata.dst_offset << 2) >> 2);
43
44    let first_vertex_or_index = src[src_base_offset + 2];
45    let vertex_or_index_count = src[src_base_offset + 0];
46
47    {
48        let can_overflow = is_bit_set(metadata.dst_offset, 30);
49        let sub_overflows = metadata.vertex_or_index_limit < first_vertex_or_index;
50        failed |= sub_overflows && !can_overflow;
51        let vertex_or_index_limit = metadata.vertex_or_index_limit - first_vertex_or_index;
52        failed |= vertex_or_index_limit < vertex_or_index_count;
53    }
54
55    let first_instance = src[src_base_offset + 3 + u32(is_indexed)];
56    let instance_count = src[src_base_offset + 1];
57
58    {
59        let can_overflow = is_bit_set(metadata.dst_offset, 31);
60        let sub_overflows = metadata.instance_limit < first_instance;
61        failed |= sub_overflows && !can_overflow;
62        let instance_limit = metadata.instance_limit - first_instance;
63        failed |= instance_limit < instance_count;
64    }
65
66    if !supports_indirect_first_instance {
67        failed |= first_instance != 0u;
68    }
69
70    let dst_offset = select(0u, 3u, write_d3d12_special_constants);
71    if failed {
72        if write_d3d12_special_constants {
73            dst[dst_base_offset + 0] = 0u;
74            dst[dst_base_offset + 1] = 0u;
75            dst[dst_base_offset + 2] = 0u;
76        }
77        dst[dst_base_offset + dst_offset + 0] = 0u;
78        dst[dst_base_offset + dst_offset + 1] = 0u;
79        dst[dst_base_offset + dst_offset + 2] = 0u;
80        dst[dst_base_offset + dst_offset + 3] = 0u;
81        if (is_indexed) {
82            dst[dst_base_offset + dst_offset + 4] = 0u;
83        }
84    } else {
85        if write_d3d12_special_constants {
86            dst[dst_base_offset + 0] = src[src_base_offset + 2 + u32(is_indexed)];
87            dst[dst_base_offset + 1] = src[src_base_offset + 3 + u32(is_indexed)];
88            dst[dst_base_offset + 2] = 0u;
89        }
90        dst[dst_base_offset + dst_offset + 0] = src[src_base_offset + 0];
91        dst[dst_base_offset + dst_offset + 1] = src[src_base_offset + 1];
92        dst[dst_base_offset + dst_offset + 2] = src[src_base_offset + 2];
93        dst[dst_base_offset + dst_offset + 3] = src[src_base_offset + 3];
94        if (is_indexed) {
95            dst[dst_base_offset + dst_offset + 4] = src[src_base_offset + 4];
96        }
97    }
98}