alex/vello
1
0
Fork 0
mirror of https://github.com/italicsjenga/vello.git synced 2025-01-10 12:41:30 +11:00
Code Issues Packages Projects Releases Wiki Activity

Merge pull request #269 from linebender/fix_fail

Browse source 
Partially revert uniform load of bump.failed
This commit is contained in:
Raph Levien 2023-01-30 16:21:35 -08:00 committed by GitHub
parent 6a184244e6 27e6fdd267
commit 872fc5c6de
No known key found for this signature in database
GPG key ID: 4AEE18F83AFDEB23
2 changed files with 6 additions and 6 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

5
shader/coarse.wgsl
View file

@ -148,15 +148,14 @@ fn main(
// Exit early if prior stages failed, as we can't run this stage. // Exit early if prior stages failed, as we can't run this stage.
// We need to check only prior stages, as if this stage has failed in another workgroup, // We need to check only prior stages, as if this stage has failed in another workgroup,
// we still want to know this workgroup's memory requirement. // we still want to know this workgroup's memory requirement.
#ifdef have_uniform
if local_id.x == 0u { if local_id.x == 0u {
// Reuse sh_part_count to hold failed flag, shmem is tight // Reuse sh_part_count to hold failed flag, shmem is tight
sh_part_count[0] = atomicLoad(&bump.failed); sh_part_count[0] = atomicLoad(&bump.failed);
} }
#ifdef have_uniform
let failed = workgroupUniformLoad(&sh_part_count[0]); let failed = workgroupUniformLoad(&sh_part_count[0]);
#else #else
workgroupBarrier(); let failed = atomicLoad(&bump.failed);
let failed = sh_part_count[0];
#endif #endif
if (failed & (STAGE_BINNING | STAGE_TILE_ALLOC | STAGE_PATH_COARSE)) != 0u { if (failed & (STAGE_BINNING | STAGE_TILE_ALLOC | STAGE_PATH_COARSE)) != 0u {
return; return;

7
shader/tile_alloc.wgsl
View file

@ -29,7 +29,9 @@ let WG_SIZE = 256u;
var<workgroup> sh_tile_count: array<u32, WG_SIZE>; var<workgroup> sh_tile_count: array<u32, WG_SIZE>;
var<workgroup> sh_tile_offset: u32; var<workgroup> sh_tile_offset: u32;
#ifdef have_uniform
var<workgroup> sh_atomic_failed: u32; var<workgroup> sh_atomic_failed: u32;
#endif
@compute @workgroup_size(256) @compute @workgroup_size(256)
fn main( fn main(
@ -39,14 +41,13 @@ fn main(
// Exit early if prior stages failed, as we can't run this stage. // Exit early if prior stages failed, as we can't run this stage.
// We need to check only prior stages, as if this stage has failed in another workgroup, // We need to check only prior stages, as if this stage has failed in another workgroup,
// we still want to know this workgroup's memory requirement. // we still want to know this workgroup's memory requirement.
#ifdef have_uniform
if local_id.x == 0u { if local_id.x == 0u {
sh_atomic_failed = atomicLoad(&bump.failed); sh_atomic_failed = atomicLoad(&bump.failed);
} }
#ifdef have_uniform
let failed = workgroupUniformLoad(&sh_atomic_failed); let failed = workgroupUniformLoad(&sh_atomic_failed);
#else #else
workgroupBarrier(); let failed = atomicLoad(&bump.failed);
let failed = sh_atomic_failed;
#endif #endif
if (failed & STAGE_BINNING) != 0u { if (failed & STAGE_BINNING) != 0u {
return; return;