| // |
| // Copyright 2019 The ANGLE Project. All rights reserved. |
| // Use of this source code is governed by a BSD-style license that can be |
| // found in the LICENSE file. |
| // |
| |
| #include "common.h" |
| |
| constant bool kCombineWithExistingResult [[function_constant(1000)]]; |
| |
| // Combine the visibility result of current render pass with previous value from previous render |
| // pass |
| struct CombineVisibilityResultOptions |
| { |
| // Start offset in the render pass's visibility buffer allocated for the query. |
| uint startOffset; |
| // How many offsets in the render pass's visibility buffer is used for the query? |
| uint numOffsets; |
| }; |
| |
| kernel void combineVisibilityResult(uint idx [[thread_position_in_grid]], |
| constant CombineVisibilityResultOptions &options [[buffer(0)]], |
| constant ushort4 *renderpassVisibilityResult [[buffer(1)]], |
| device ushort4 *finalResults [[buffer(2)]]) |
| { |
| if (idx > 0) |
| { |
| // NOTE(hqle): |
| // This is a bit wasteful to use a WARP of multiple threads just for combining one integer. |
| // Consider a better approach. |
| return; |
| } |
| ushort4 finalResult16x4; |
| |
| if (kCombineWithExistingResult) |
| { |
| finalResult16x4 = finalResults[0]; |
| } |
| else |
| { |
| finalResult16x4 = ushort4(0, 0, 0, 0); |
| } |
| |
| for (uint i = 0; i < options.numOffsets; ++i) |
| { |
| uint offset = options.startOffset + i; |
| ushort4 renderpassResult = renderpassVisibilityResult[offset]; |
| |
| // Only boolean result is required, so bitwise OR is enough |
| finalResult16x4 = finalResult16x4 | renderpassResult; |
| } |
| finalResults[0] = finalResult16x4; |
| } |