From 137753ae28df04f806d9de9c95ab05592c0264ae Mon Sep 17 00:00:00 2001 From: Julia Tatz Date: Thu, 25 May 2023 20:28:45 -0400 Subject: [PATCH] zink: fix layout(local_size_variable) for vk1.3+ 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: (cherry picked from commit 2f3841339cd85f1953be2613688b86e7372e946c) --- .pick_status.json | 2 +- src/gallium/drivers/zink/nir_to_spirv/nir_to_spirv.c | 12 ++++++------ .../drivers/zink/nir_to_spirv/spirv_builder.c | 12 ++++++++++++ .../drivers/zink/nir_to_spirv/spirv_builder.h | 3 +++ 4 files changed, 22 insertions(+), 7 deletions(-) diff --git a/.pick_status.json b/.pick_status.json index cc59e580baa..300aed8f379 100644 --- a/.pick_status.json +++ b/.pick_status.json @@ -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" }, diff --git a/src/gallium/drivers/zink/nir_to_spirv/nir_to_spirv.c b/src/gallium/drivers/zink/nir_to_spirv/nir_to_spirv.c index 32f0934be41..ff969b9b8fd 100644 --- a/src/gallium/drivers/zink/nir_to_spirv/nir_to_spirv.c +++ b/src/gallium/drivers/zink/nir_to_spirv/nir_to_spirv.c @@ -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); } } diff --git a/src/gallium/drivers/zink/nir_to_spirv/spirv_builder.c b/src/gallium/drivers/zink/nir_to_spirv/spirv_builder.c index 557fbfa6001..7135a1ce69e 100644 --- a/src/gallium/drivers/zink/nir_to_spirv/spirv_builder.c +++ b/src/gallium/drivers/zink/nir_to_spirv/spirv_builder.c @@ -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) diff --git a/src/gallium/drivers/zink/nir_to_spirv/spirv_builder.h b/src/gallium/drivers/zink/nir_to_spirv/spirv_builder.h index 10d2361ee16..9286bba12b2 100644 --- a/src/gallium/drivers/zink/nir_to_spirv/spirv_builder.h +++ b/src/gallium/drivers/zink/nir_to_spirv/spirv_builder.h @@ -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);