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

zink: fix layout(local_size_variable) for vk1.3+

Browse source 
Use the correct exec-mode op for LocalSizeId
Corrected typo `gl_LocalGroupSize` -> `gl_LocalGroupSizeARB`

Fixes: 99bd1eaf ("zink: use spir-v 1.6 local-size when needed")
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/23244>
(cherry picked from commit 2f3841339c)
This commit is contained in:
Julia Tatz 2023-05-25 20:28:45 -04:00 • committed by Eric Engestrom
parent cdbe5cef34
commit 137753ae28
4 changed files with 22 additions and 7 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

2
.pick_status.json
View file

@ -643,7 +643,7 @@
"description": "zink: fix layout(local_size_variable) for vk1.3+",
"nominated": true,
"nomination_type": 1,
"resolution": 0,
"resolution": 1,
"main_sha": null,
"because_sha": "99bd1eaf3d20184abaff91e63d8aacded79d4d74"
},

12
src/gallium/drivers/zink/nir_to_spirv/nir_to_spirv.c
View file

@ -4697,17 +4697,17 @@ nir_to_spirv(struct nir_shader *s, const struct zink_shader_info *sinfo, uint32_
spirv_builder_emit_specid(&ctx.builder, sizes[i], ids[i]);
spirv_builder_emit_name(&ctx.builder, sizes[i], names[i]);
}
SpvId var_type = get_uvec_type(&ctx, 32, 3);
// Even when using LocalSizeId this need to be initialized for nir_intrinsic_load_workgroup_size
ctx.local_group_size_var = spirv_builder_spec_const_composite(&ctx.builder, var_type, sizes, 3);
spirv_builder_emit_name(&ctx.builder, ctx.local_group_size_var, "gl_LocalGroupSizeARB");
/* WorkgroupSize is deprecated in SPIR-V 1.6 */
if (spirv_version >= SPIRV_VERSION(1, 6)) {
uint32_t sizes32[] = { sizes[0], sizes[1], sizes[2] };
spirv_builder_emit_exec_mode_literal3(&ctx.builder, entry_point,
spirv_builder_emit_exec_mode_id3(&ctx.builder, entry_point,
SpvExecutionModeLocalSizeId,
sizes32);
sizes);
} else {
SpvId var_type = get_uvec_type(&ctx, 32, 3);
ctx.local_group_size_var = spirv_builder_spec_const_composite(&ctx.builder, var_type, sizes, 3);
spirv_builder_emit_name(&ctx.builder, ctx.local_group_size_var, "gl_LocalGroupSize");
spirv_builder_emit_builtin(&ctx.builder, ctx.local_group_size_var, SpvBuiltInWorkgroupSize);
}
}

12
src/gallium/drivers/zink/nir_to_spirv/spirv_builder.c
View file

@ -173,6 +173,18 @@ spirv_builder_emit_exec_mode_literal3(struct spirv_builder *b, SpvId entry_point
spirv_buffer_emit_word(&b->exec_modes, param[i]);
}
void
spirv_builder_emit_exec_mode_id3(struct spirv_builder *b, SpvId entry_point,
SpvExecutionMode exec_mode, SpvId param[3])
{
spirv_buffer_prepare(&b->exec_modes, b->mem_ctx, 6);
spirv_buffer_emit_word(&b->exec_modes, SpvOpExecutionModeId | (6 << 16));
spirv_buffer_emit_word(&b->exec_modes, entry_point);
spirv_buffer_emit_word(&b->exec_modes, exec_mode);
for (unsigned i = 0; i < 3; i++)
spirv_buffer_emit_word(&b->exec_modes, param[i]);
}
void
spirv_builder_emit_exec_mode(struct spirv_builder *b, SpvId entry_point,
SpvExecutionMode exec_mode)

3
src/gallium/drivers/zink/nir_to_spirv/spirv_builder.h
View file

@ -157,6 +157,9 @@ void
spirv_builder_emit_exec_mode_literal3(struct spirv_builder *b, SpvId entry_point,
SpvExecutionMode exec_mode, uint32_t param[3]);
void
spirv_builder_emit_exec_mode_id3(struct spirv_builder *b, SpvId entry_point,
SpvExecutionMode exec_mode, SpvId param[3]);
void
spirv_builder_emit_exec_mode(struct spirv_builder *b, SpvId entry_point,
SpvExecutionMode exec_mode);