fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2026-06-17 02:18:21 +02:00
Code Issues Projects Releases Packages Wiki Activity Actions 2

radv: fix local invocation index for mesh/task and quad derivatives on GFX12

Browse source 
It must be lowered.

This fixes
dEQP-VK.spirv_assembly.instruction.compute.compute_shader_derivatives.{mesh,task}.*.

Cc: mesa-stable
Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40194>
This commit is contained in:
Samuel Pitoiset 2026-03-03 10:18:50 +01:00 • committed by Marge Bot
parent 4428541c54
commit 3c4cb16159
1 changed files with 13 additions and 4 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

17
src/amd/vulkan/radv_shader.c
View file

@ -662,15 +662,24 @@ radv_shader_spirv_to_nir(struct radv_device *device, const struct radv_shader_st
NIR_PASS(_, nir, nir_lower_compute_system_values, &csv_options);
}
bool lower_local_invocation_index = false;
if (nir->info.derivative_group == DERIVATIVE_GROUP_QUADS &&
((nir->info.stage == MESA_SHADER_COMPUTE || nir->info.stage == MESA_SHADER_TASK ||
(nir->info.stage == MESA_SHADER_MESH && pdev->info.mesh_fast_launch_2)))) {
lower_local_invocation_index = true;
} else if (nir->info.stage == MESA_SHADER_COMPUTE &&
(((nir->info.workgroup_size[0] == 1) + (nir->info.workgroup_size[1] == 1) +
(nir->info.workgroup_size[2] == 1)) == 2)) {
lower_local_invocation_index = true;
}
nir_lower_compute_system_values_options csv_options = {
/* Mesh shaders run as NGG which can implement local_invocation_index from
* the wave ID in merged_wave_info, but they don't have local_invocation_ids on GFX10.3.
*/
.lower_cs_local_id_to_index = nir->info.stage == MESA_SHADER_MESH && !pdev->info.mesh_fast_launch_2,
.lower_local_invocation_index = nir->info.stage == MESA_SHADER_COMPUTE &&
((((nir->info.workgroup_size[0] == 1) + (nir->info.workgroup_size[1] == 1) +
(nir->info.workgroup_size[2] == 1)) == 2) ||
nir->info.derivative_group == DERIVATIVE_GROUP_QUADS),
.lower_local_invocation_index = lower_local_invocation_index,
};
NIR_PASS(_, nir, nir_lower_compute_system_values, &csv_options);