From: Jesse Natalie Date: Fri, 31 Mar 2023 17:25:00 +0000 (-0700) Subject: microsoft/compiler: Assign 1D wave IDs based on local thread ID X-Git-Tag: upstream/23.3.3~10385 X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=4f56cede6deff561a8cfa3fe893cf085cb8b2f1c;p=platform%2Fupstream%2Fmesa.git microsoft/compiler: Assign 1D wave IDs based on local thread ID Fixes corruption/flickering seen in DOOM Eternal's decals/lighting. It seems the shader has an implicit assumption about this property. Part-of: --- diff --git a/src/microsoft/compiler/dxil_nir.c b/src/microsoft/compiler/dxil_nir.c index fa2eb35..6530914 100644 --- a/src/microsoft/compiler/dxil_nir.c +++ b/src/microsoft/compiler/dxil_nir.c @@ -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);