1// 2// Copyright 2019 The ANGLE Project. All rights reserved. 3// Use of this source code is governed by a BSD-style license that can be 4// found in the LICENSE file. 5// 6 7#include "common.h" 8 9constant bool kCombineWithExistingResult [[function_constant(1000)]]; 10 11// Combine the visibility result of current render pass with previous value from previous render 12// pass 13struct CombineVisibilityResultOptions 14{ 15 // Start offset in the render pass's visibility buffer allocated for the query. 16 uint startOffset; 17 // How many offsets in the render pass's visibility buffer is used for the query? 18 uint numOffsets; 19}; 20 21kernel void combineVisibilityResult(uint idx [[thread_position_in_grid]], 22 constant CombineVisibilityResultOptions &options [[buffer(0)]], 23 constant ushort4 *renderpassVisibilityResult [[buffer(1)]], 24 device ushort4 *finalResults [[buffer(2)]]) 25{ 26 if (idx > 0) 27 { 28 // NOTE(hqle): 29 // This is a bit wasteful to use a WARP of multiple threads just for combining one integer. 30 // Consider a better approach. 31 return; 32 } 33 ushort4 finalResult16x4; 34 35 if (kCombineWithExistingResult) 36 { 37 finalResult16x4 = finalResults[0]; 38 } 39 else 40 { 41 finalResult16x4 = ushort4(0, 0, 0, 0); 42 } 43 44 for (uint i = 0; i < options.numOffsets; ++i) 45 { 46 uint offset = options.startOffset + i; 47 ushort4 renderpassResult = renderpassVisibilityResult[offset]; 48 49 // Only boolean result is required, so bitwise OR is enough 50 finalResult16x4 = finalResult16x4 | renderpassResult; 51 } 52 finalResults[0] = finalResult16x4; 53} 54