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