fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2026-01-01 05:20:09 +01:00
Code Issues Projects Releases Packages Wiki Activity Actions 1

microsoft/compiler: Assign 1D wave IDs based on local thread ID

Browse source 
Fixes corruption/flickering seen in DOOM Eternal's decals/lighting.
It seems the shader has an implicit assumption about this property.

Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/22225>
This commit is contained in:
Jesse Natalie 2023-03-31 10:25:00 -07:00 • committed by Marge Bot
parent eeb67362da
commit 4f56cede6d
1 changed files with 10 additions and 1 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

11
src/microsoft/compiler/dxil_nir.c
View file

@ -2061,11 +2061,20 @@ lower_subgroup_id(nir_builder *b, nir_instr *instr, void *data)
if (intr->intrinsic != nir_intrinsic_load_subgroup_id)
return false;
b->cursor = nir_before_block(nir_start_block(b->impl));
if (b->shader->info.workgroup_size[1] == 1 &&
b->shader->info.workgroup_size[2] == 1) {
/* When using Nx1x1 groups, use a simple stable algorithm
* which is almost guaranteed to be correct. */
nir_ssa_def *subgroup_id = nir_udiv(b, nir_load_local_invocation_index(b), nir_load_subgroup_size(b));
nir_ssa_def_rewrite_uses(&intr->dest.ssa, subgroup_id);
return true;
}
nir_ssa_def **subgroup_id = (nir_ssa_def **)data;
if (*subgroup_id == NULL) {
nir_variable *subgroup_id_counter = nir_variable_create(b->shader, nir_var_mem_shared, glsl_uint_type(), "dxil_SubgroupID_counter");
nir_variable *subgroup_id_local = nir_local_variable_create(b->impl, glsl_uint_type(), "dxil_SubgroupID_local");
b->cursor = nir_before_block(nir_start_block(b->impl));
nir_store_var(b, subgroup_id_local, nir_imm_int(b, 0), 1);
nir_deref_instr *counter_deref = nir_build_deref_var(b, subgroup_id_counter);