From 04d3b3bde5835e156c23e8e602be70494ab8b14f Mon Sep 17 00:00:00 2001 From: Georg Lehmann Date: Tue, 9 Sep 2025 18:48:27 +0200 Subject: [PATCH] mesa,glsl,spirv: set new subgroup size info MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Reviewed-by: Marek Olšák Acked-by: Timur Kristóf Part-of: --- src/compiler/glsl/gl_nir_linker.c | 6 ++++++ src/compiler/spirv/nir_spirv.h | 2 ++ src/compiler/spirv/spirv_to_nir.c | 8 ++++++++ src/gallium/auxiliary/nir/tgsi_to_nir.c | 2 ++ src/mesa/main/ff_fragment_shader.c | 1 + src/mesa/main/glspirv.c | 4 ++++ 6 files changed, 23 insertions(+) diff --git a/src/compiler/glsl/gl_nir_linker.c b/src/compiler/glsl/gl_nir_linker.c index 82b2bd1cdf2..a585b88fd2f 100644 --- a/src/compiler/glsl/gl_nir_linker.c +++ b/src/compiler/glsl/gl_nir_linker.c @@ -2697,6 +2697,12 @@ link_intrastage_shaders(void *mem_ctx, link_layer_viewport_relative_qualifier(prog, gl_prog, shader_list, num_shaders); gl_prog->nir->info.view_mask = view_mask; + gl_prog->nir->info.api_subgroup_size_draw_uniform = + !mesa_shader_stage_uses_workgroup(gl_prog->nir->info.stage); + if (KHR_shader_subgroup_basic_enable) { + gl_prog->nir->info.api_subgroup_size = ctx->screen->caps.shader_subgroup_size; + gl_prog->nir->info.max_subgroup_size = ctx->screen->caps.shader_subgroup_size; + } gl_prog->nir->info.subgroup_size = KHR_shader_subgroup_basic_enable ? SUBGROUP_SIZE_API_CONSTANT : SUBGROUP_SIZE_UNIFORM; diff --git a/src/compiler/spirv/nir_spirv.h b/src/compiler/spirv/nir_spirv.h index 6177ad892f0..235c064bf8b 100644 --- a/src/compiler/spirv/nir_spirv.h +++ b/src/compiler/spirv/nir_spirv.h @@ -152,6 +152,8 @@ struct spirv_to_nir_options { /* Shader index provided by VkPipelineShaderStageNodeCreateInfoAMDX */ uint32_t shader_index; + /* If GroupNonUniform capability is used, set this api subgroup size. */ + uint8_t group_non_uniform_subgroup_size; }; enum spirv_verify_result { diff --git a/src/compiler/spirv/spirv_to_nir.c b/src/compiler/spirv/spirv_to_nir.c index 2377c410c82..8e16988994e 100644 --- a/src/compiler/spirv/spirv_to_nir.c +++ b/src/compiler/spirv/spirv_to_nir.c @@ -5659,6 +5659,9 @@ vtn_handle_execution_mode(struct vtn_builder *b, struct vtn_value *entry_point, vtn_assert(b->shader->info.stage == MESA_SHADER_KERNEL); vtn_assert(b->shader->info.subgroup_size == SUBGROUP_SIZE_VARYING); b->shader->info.subgroup_size = mode->operands[0]; + b->shader->info.api_subgroup_size = mode->operands[0]; + b->shader->info.max_subgroup_size = mode->operands[0]; + b->shader->info.min_subgroup_size = mode->operands[0]; break; case SpvExecutionModeSubgroupsPerWorkgroup: @@ -7232,6 +7235,11 @@ spirv_to_nir(const uint32_t *words, size_t word_count, b->enabled_capabilities.GroupNonUniform) b->shader->info.subgroup_size = SUBGROUP_SIZE_API_CONSTANT; + if (b->enabled_capabilities.GroupNonUniform && options->group_non_uniform_subgroup_size) { + b->shader->info.api_subgroup_size = options->group_non_uniform_subgroup_size; + b->shader->info.max_subgroup_size = options->group_non_uniform_subgroup_size; + } + /* DirectXShaderCompiler and glslang/shaderc both create OpKill from HLSL's * discard/clip, which uses demote semantics. DirectXShaderCompiler will use * demote if the extension is enabled, so we disable this workaround in that diff --git a/src/gallium/auxiliary/nir/tgsi_to_nir.c b/src/gallium/auxiliary/nir/tgsi_to_nir.c index a671148a3ae..165c97619f9 100644 --- a/src/gallium/auxiliary/nir/tgsi_to_nir.c +++ b/src/gallium/auxiliary/nir/tgsi_to_nir.c @@ -2352,6 +2352,8 @@ ttn_compile_init(const void *tgsi_tokens, } } + s->info.api_subgroup_size_draw_uniform = s->info.stage != MESA_SHADER_COMPUTE; + if (s->info.stage == MESA_SHADER_COMPUTE && (!s->info.workgroup_size[0] || !s->info.workgroup_size[1] || diff --git a/src/mesa/main/ff_fragment_shader.c b/src/mesa/main/ff_fragment_shader.c index 235ab7e0aaa..5a1ffa2f4ff 100644 --- a/src/mesa/main/ff_fragment_shader.c +++ b/src/mesa/main/ff_fragment_shader.c @@ -945,6 +945,7 @@ create_new_program(struct gl_context *ctx, struct state_key *key, s->info.separate_shader = true; s->info.subgroup_size = SUBGROUP_SIZE_UNIFORM; + s->info.api_subgroup_size_draw_uniform = true; s->info.io_lowered = true; p.b = &b; diff --git a/src/mesa/main/glspirv.c b/src/mesa/main/glspirv.c index 3213e9a04b6..d8628268eb1 100644 --- a/src/mesa/main/glspirv.c +++ b/src/mesa/main/glspirv.c @@ -36,6 +36,8 @@ #include "util/u_atomic.h" #include "api_exec_decl.h" +#include "pipe/p_screen.h" + void _mesa_spirv_module_reference(struct gl_spirv_module **dest, struct gl_spirv_module *src) @@ -274,6 +276,7 @@ _mesa_spirv_to_nir(struct gl_context *ctx, */ .shared_addr_format = nir_address_format_32bit_offset, + .group_non_uniform_subgroup_size = ctx->screen->caps.shader_subgroup_size, }; nir_shader *nir = @@ -297,6 +300,7 @@ _mesa_spirv_to_nir(struct gl_context *ctx, nir_validate_shader(nir, "after spirv_to_nir"); nir->info.separate_shader = linked_shader->Program->info.separate_shader; + nir->info.api_subgroup_size_draw_uniform = !mesa_shader_stage_uses_workgroup(stage); /* Convert some sysvals to input varyings. */ const struct nir_lower_sysvals_to_varyings_options sysvals_to_varyings = {