ata[dupIndex]; outputParams.data[outIndex + 1u] = inputParams.data[dupIndex + 1u]; outIndex = outIndex + 2u; } for(var i = 0u; i < numInputParams; i = i + 1u) { outputParams.data[outIndex + i] = inputParams.data[inIndex + i]; } } @compute @workgroup_size(64, 1, 1) fn main(@builtin(global_invocation_id) id : vec3) { if (id.x >= batch.numDraws) { return; } if(!bool(batch.flags & kValidationEnabled)) { pass(id.x); return; } let inputIndex = batch.indirectOffsets[id.x]; if(!bool(batch.flags & kIndirectFirstInstanceEnabled)) { // firstInstance is always the last parameter let firstInstance = inputParams.data[inputIndex + numIndirectParamsPerDrawCallInput() - 1u]; if (firstInstance != 0u) { fail(id.x); return; } } if (!bool(batch.flags & kIndexedDraw)) { pass(id.x); return; } if (batch.numIndexBufferElementsHigh >= 2u) { // firstIndex and indexCount are both u32. The maximum possible sum of these // values is 0x1fffffffe, which is less than 0x200000000. Nothing to validate. pass(id.x); return; } let firstIndex = inputParams.data[inputIndex + kFirstIndexEntry]; if (batch.numIndexBufferElementsHigh == 0u && batch.numIndexBufferElementsLow < firstIndex) { fail(id.x); return; } // Note that this subtraction may underflow, but only when // numIndexBufferElementsHigh is 1u. The result is still correct in that case. let maxIndexCount = batch.numIndexBufferElementsLow - firstIndex; let indexCount = inputParams.data[inputIndex + kIndexCountEntry]; if (indexCount > maxIndexCount) { fail(id.x); return; } pass(id.x); }