fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2026-05-06 20:18:12 +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>
(cherry picked from commit 3c4cb16159)

Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40359>
This commit is contained in:
Samuel Pitoiset 2026-03-03 10:18:50 +01:00 • committed by Eric Engestrom
parent f858d2238e
commit ecb7bf7b68
2 changed files with 14 additions and 5 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

2
.pick_status.json
View file

@ -1304,7 +1304,7 @@
"description": "radv: fix local invocation index for mesh/task and quad derivatives on GFX12",
"nominated": true,
"nomination_type": 1,
"resolution": 0,
"resolution": 1,
"main_sha": null,
"because_sha": null,
"notes": null

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

@ -655,15 +655,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);