From ce91b0be08319556e23a3366c0ff18ff34ac0775 Mon Sep 17 00:00:00 2001 From: Georg Lehmann Date: Tue, 9 Sep 2025 16:41:00 +0200 Subject: [PATCH] nir: define new subgroup size info MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Reviewed-by: Marek Olšák Reviewed-by: Timur Kristóf Part-of: --- src/compiler/nir/nir.c | 2 ++ src/compiler/nir/nir_print.c | 5 +++ src/compiler/nir/nir_validate.c | 13 ++++++++ .../nir/tests/lower_discard_if_tests.cpp | 11 ++++--- .../tests/minimize_call_live_states_test.cpp | 15 ++++++--- src/compiler/nir/tests/opt_loop_tests.cpp | 27 ++++++++++------ src/compiler/shader_info.h | 31 +++++++++++++++++-- src/mesa/program/prog_to_nir.c | 2 ++ src/mesa/state_tracker/st_atifs_to_nir.c | 2 ++ 9 files changed, 88 insertions(+), 20 deletions(-) diff --git a/src/compiler/nir/nir.c b/src/compiler/nir/nir.c index b5a3f9e9d36..0a2d9fae0db 100644 --- a/src/compiler/nir/nir.c +++ b/src/compiler/nir/nir.c @@ -217,6 +217,8 @@ nir_shader_create(void *mem_ctx, */ shader->info.prev_stage = MESA_SHADER_NONE; shader->info.next_stage = MESA_SHADER_NONE; + shader->info.max_subgroup_size = 128; + shader->info.min_subgroup_size = 1; exec_list_make_empty(&shader->functions); diff --git a/src/compiler/nir/nir_print.c b/src/compiler/nir/nir_print.c index 3fb633fad46..a7b6e7a85b3 100644 --- a/src/compiler/nir/nir_print.c +++ b/src/compiler/nir/nir_print.c @@ -2727,6 +2727,11 @@ print_shader_info(const struct shader_info *info, FILE *fp) fprintf(fp, "subgroup_size: %u\n", info->subgroup_size); + print_nz_unsigned(fp, "api_subgroup_size", info->api_subgroup_size); + fprintf(fp, "max_subgroup_size: %u\n", info->max_subgroup_size); + fprintf(fp, "min_subgroup_size: %u\n", info->min_subgroup_size); + + print_nz_bool(fp, "api_subgroup_size_draw_uniform", info->api_subgroup_size_draw_uniform); print_nz_bool(fp, "uses_wide_subgroup_intrinsics", info->uses_wide_subgroup_intrinsics); bool has_xfb_stride = info->xfb_stride[0] || info->xfb_stride[1] || info->xfb_stride[2] || info->xfb_stride[3]; diff --git a/src/compiler/nir/nir_validate.c b/src/compiler/nir/nir_validate.c index 2f724ae3748..5f40a83d3eb 100644 --- a/src/compiler/nir/nir_validate.c +++ b/src/compiler/nir/nir_validate.c @@ -2225,6 +2225,19 @@ nir_validate_shader(nir_shader *shader, const char *when) validate_assert(&state, shader->xfb_info->output_count > 0); } + validate_assert(&state, + util_is_power_of_two_or_zero(shader->info.api_subgroup_size)); + validate_assert(&state, + util_is_power_of_two_nonzero(shader->info.max_subgroup_size)); + validate_assert(&state, + util_is_power_of_two_nonzero(shader->info.min_subgroup_size)); + validate_assert(&state, shader->info.max_subgroup_size <= 128); + validate_assert(&state, shader->info.min_subgroup_size <= 128); + validate_assert(&state, shader->info.api_subgroup_size == 0 || + shader->info.api_subgroup_size >= shader->info.max_subgroup_size); + validate_assert(&state, + shader->info.min_subgroup_size <= shader->info.max_subgroup_size); + if (_mesa_hash_table_num_entries(state.errors) > 0) dump_errors(&state, when); diff --git a/src/compiler/nir/tests/lower_discard_if_tests.cpp b/src/compiler/nir/tests/lower_discard_if_tests.cpp index e43dfcca663..00a06910763 100644 --- a/src/compiler/nir/tests/lower_discard_if_tests.cpp +++ b/src/compiler/nir/tests/lower_discard_if_tests.cpp @@ -36,7 +36,8 @@ TEST_F(nir_lower_discard_if_test, move_single_terminate_out_of_loop) check_nir_string(NIR_REFERENCE_SHADER(R"( shader: MESA_SHADER_FRAGMENT name: nir_lower_discard_if_test - subgroup_size: 0 + max_subgroup_size: 128 + min_subgroup_size: 1 decl_var shader_in INTERP_MODE_SMOOTH none int in (VARYING_SLOT_POS.x, 0, 0) decl_function main () (entrypoint) @@ -94,7 +95,8 @@ TEST_F(nir_lower_discard_if_test, move_multiple_terminate_out_of_loop) check_nir_string(NIR_REFERENCE_SHADER(R"( shader: MESA_SHADER_FRAGMENT name: nir_lower_discard_if_test - subgroup_size: 0 + max_subgroup_size: 128 + min_subgroup_size: 1 decl_var shader_in INTERP_MODE_SMOOTH none int in (VARYING_SLOT_POS.x, 0, 0) decl_function main () (entrypoint) @@ -165,7 +167,8 @@ TEST_F(nir_lower_discard_if_test, move_terminate_out_of_nested_loop) check_nir_string(NIR_REFERENCE_SHADER(R"( shader: MESA_SHADER_FRAGMENT name: nir_lower_discard_if_test - subgroup_size: 0 + max_subgroup_size: 128 + min_subgroup_size: 1 decl_var shader_in INTERP_MODE_SMOOTH none int in (VARYING_SLOT_POS.x, 0, 0) decl_function main () (entrypoint) @@ -225,4 +228,4 @@ TEST_F(nir_lower_discard_if_test, move_terminate_out_of_nested_loop) block b14: } )")); -} \ No newline at end of file +} diff --git a/src/compiler/nir/tests/minimize_call_live_states_test.cpp b/src/compiler/nir/tests/minimize_call_live_states_test.cpp index 5bf3afef21b..4aac6869bde 100644 --- a/src/compiler/nir/tests/minimize_call_live_states_test.cpp +++ b/src/compiler/nir/tests/minimize_call_live_states_test.cpp @@ -39,7 +39,8 @@ TEST_F(nir_minimize_call_live_states_test, no_live_states) shader: MESA_SHADER_COMPUTE name: nir_minimize_call_live_states_test workgroup_size: 1, 1, 1 - subgroup_size: 0 + max_subgroup_size: 128 + min_subgroup_size: 1 decl_function main () (entrypoint) impl main { @@ -76,7 +77,8 @@ TEST_F(nir_minimize_call_live_states_test, life_intrinsics) shader: MESA_SHADER_COMPUTE name: nir_minimize_call_live_states_test workgroup_size: 1, 1, 1 - subgroup_size: 0 + max_subgroup_size: 128 + min_subgroup_size: 1 decl_function main () (entrypoint) impl main { @@ -123,7 +125,8 @@ TEST_F(nir_minimize_call_live_states_test, life_alu) shader: MESA_SHADER_COMPUTE name: nir_minimize_call_live_states_test workgroup_size: 1, 1, 1 - subgroup_size: 0 + max_subgroup_size: 128 + min_subgroup_size: 1 decl_function main () (entrypoint) impl main { @@ -174,7 +177,8 @@ TEST_F(nir_minimize_call_live_states_test, call_inside_if) shader: MESA_SHADER_COMPUTE name: nir_minimize_call_live_states_test workgroup_size: 1, 1, 1 - subgroup_size: 0 + max_subgroup_size: 128 + min_subgroup_size: 1 decl_function main () (entrypoint) impl main { @@ -249,7 +253,8 @@ TEST_F(nir_minimize_call_live_states_test, call_inside_loop) shader: MESA_SHADER_COMPUTE name: nir_minimize_call_live_states_test workgroup_size: 1, 1, 1 - subgroup_size: 0 + max_subgroup_size: 128 + min_subgroup_size: 1 decl_function main () (entrypoint) impl main { diff --git a/src/compiler/nir/tests/opt_loop_tests.cpp b/src/compiler/nir/tests/opt_loop_tests.cpp index dbc5026af6f..affa612d3c7 100644 --- a/src/compiler/nir/tests/opt_loop_tests.cpp +++ b/src/compiler/nir/tests/opt_loop_tests.cpp @@ -136,7 +136,8 @@ TEST_F(nir_opt_loop_test, opt_loop_merge_terminators_break_in_then) check_nir_string(NIR_REFERENCE_SHADER(R"( shader: MESA_SHADER_FRAGMENT name: nir_opt_loop_test - subgroup_size: 0 + max_subgroup_size: 128 + min_subgroup_size: 1 decl_var shader_in INTERP_MODE_SMOOTH none int in (VARYING_SLOT_POS.x, 0, 0) decl_var shader_out INTERP_MODE_NONE none int out (FRAG_RESULT_DEPTH.x, 0, 0) decl_var ubo INTERP_MODE_NONE none int ubo1 (0, 0, 0) @@ -191,7 +192,8 @@ TEST_F(nir_opt_loop_test, opt_loop_merge_terminators_break_in_else) check_nir_string(NIR_REFERENCE_SHADER(R"( shader: MESA_SHADER_FRAGMENT name: nir_opt_loop_test - subgroup_size: 0 + max_subgroup_size: 128 + min_subgroup_size: 1 decl_var shader_in INTERP_MODE_SMOOTH none int in (VARYING_SLOT_POS.x, 0, 0) decl_var shader_out INTERP_MODE_NONE none int out (FRAG_RESULT_DEPTH.x, 0, 0) decl_var ubo INTERP_MODE_NONE none int ubo1 (0, 0, 0) @@ -264,7 +266,8 @@ TEST_F(nir_opt_loop_test, opt_loop_merge_terminators_deref_after_first_if) check_nir_string(NIR_REFERENCE_SHADER(R"( shader: MESA_SHADER_FRAGMENT name: nir_opt_loop_test - subgroup_size: 0 + max_subgroup_size: 128 + min_subgroup_size: 1 decl_var shader_in INTERP_MODE_SMOOTH none int in (VARYING_SLOT_POS.x, 0, 0) decl_var shader_out INTERP_MODE_NONE none int out (FRAG_RESULT_DEPTH.x, 0, 0) decl_var ubo INTERP_MODE_NONE none int ubo1 (0, 0, 0) @@ -343,7 +346,8 @@ TEST_F(nir_opt_loop_test, opt_loop_merge_terminators_deref_phi_index) check_nir_string(NIR_REFERENCE_SHADER(R"( shader: MESA_SHADER_FRAGMENT name: nir_opt_loop_test - subgroup_size: 0 + max_subgroup_size: 128 + min_subgroup_size: 1 decl_var shader_in INTERP_MODE_SMOOTH none int in (VARYING_SLOT_POS.x, 0, 0) decl_var shader_out INTERP_MODE_NONE none int out (FRAG_RESULT_DEPTH.x, 0, 0) decl_var ubo INTERP_MODE_NONE none int ubo1 (0, 0, 0) @@ -427,7 +431,8 @@ TEST_F(nir_opt_loop_test, opt_loop_merge_terminators_skip_merge_if_phis) check_nir_string(NIR_REFERENCE_SHADER(R"( shader: MESA_SHADER_FRAGMENT name: nir_opt_loop_test - subgroup_size: 0 + max_subgroup_size: 128 + min_subgroup_size: 1 decl_var shader_in INTERP_MODE_SMOOTH none int in (VARYING_SLOT_POS.x, 0, 0) decl_var shader_out INTERP_MODE_NONE none int out (FRAG_RESULT_DEPTH.x, 0, 0) decl_var ubo INTERP_MODE_NONE none int ubo1 (0, 0, 0) @@ -506,7 +511,8 @@ TEST_F(nir_opt_loop_test, opt_loop_merge_terminators_skip_merge_if_phis_nested_l check_nir_string(NIR_REFERENCE_SHADER(R"( shader: MESA_SHADER_FRAGMENT name: nir_opt_loop_test - subgroup_size: 0 + max_subgroup_size: 128 + min_subgroup_size: 1 decl_var shader_in INTERP_MODE_SMOOTH none int in (VARYING_SLOT_POS.x, 0, 0) decl_var shader_out INTERP_MODE_NONE none int out (FRAG_RESULT_DEPTH.x, 0, 0) decl_var ubo INTERP_MODE_NONE none int ubo1 (0, 0, 0) @@ -581,7 +587,8 @@ TEST_F(nir_opt_loop_test, opt_loop_peel_initial_break_ends_with_jump) check_nir_string(NIR_REFERENCE_SHADER(R"( shader: MESA_SHADER_FRAGMENT name: nir_opt_loop_test - subgroup_size: 0 + max_subgroup_size: 128 + min_subgroup_size: 1 decl_var shader_in INTERP_MODE_SMOOTH none int in (VARYING_SLOT_POS.x, 0, 0) decl_var shader_out INTERP_MODE_NONE none int out (FRAG_RESULT_DEPTH.x, 0, 0) decl_var ubo INTERP_MODE_NONE none int ubo1 (0, 0, 0) @@ -652,7 +659,8 @@ TEST_F(nir_opt_loop_test, opt_loop_peel_initial_break_nontrivial_break) check_nir_string(NIR_REFERENCE_SHADER(R"( shader: MESA_SHADER_FRAGMENT name: nir_opt_loop_test - subgroup_size: 0 + max_subgroup_size: 128 + min_subgroup_size: 1 decl_var shader_in INTERP_MODE_SMOOTH none int in (VARYING_SLOT_POS.x, 0, 0) decl_var shader_out INTERP_MODE_NONE none int out (FRAG_RESULT_DEPTH.x, 0, 0) decl_var ubo INTERP_MODE_NONE none int ubo1 (0, 0, 0) @@ -725,7 +733,8 @@ TEST_F(nir_opt_loop_test, opt_loop_peel_initial_break_deref) check_nir_string(NIR_REFERENCE_SHADER(R"( shader: MESA_SHADER_FRAGMENT name: nir_opt_loop_test - subgroup_size: 0 + max_subgroup_size: 128 + min_subgroup_size: 1 decl_var shader_in INTERP_MODE_SMOOTH none int in (VARYING_SLOT_POS.x, 0, 0) decl_var shader_out INTERP_MODE_NONE none int out (FRAG_RESULT_DEPTH.x, 0, 0) decl_var ubo INTERP_MODE_NONE none int ubo1 (0, 0, 0) diff --git a/src/compiler/shader_info.h b/src/compiler/shader_info.h index ccac90e0354..772dc1e20b4 100644 --- a/src/compiler/shader_info.h +++ b/src/compiler/shader_info.h @@ -172,13 +172,40 @@ typedef struct shader_info { */ uint16_t workgroup_size[3]; - enum gl_subgroup_size subgroup_size; uint8_t num_subgroups; + enum gl_subgroup_size subgroup_size; /* To be removed. */ + + /* The value reported in gl_SubgroupSize. + * Must be a power of two between 1 and 128 + * or 0 if still unknown. + */ + uint8_t api_subgroup_size; + + /* The maximum subgroup size dispatched by the hw. + * Must be a power of two between 1 and 128. + * Must not be larger than api_subgroup_size, + * (unless api_subgroup_size is 0). + */ + uint8_t max_subgroup_size; + + /* The minimum subgroup size dispatched by the hw. + * Must be a power of two between 1 and 128. + * Must not be larger than max_subgroup_size. + */ + uint8_t min_subgroup_size; + + /* api_subgroup_size must appear to be uniform in + * the current stage for a whole draw. + * There is no equivalent for dispatches, + * because it would be required to be true. + */ + bool api_subgroup_size_draw_uniform:1; + /** * Uses subgroup intrinsics which can communicate across a quad. */ - bool uses_wide_subgroup_intrinsics; + bool uses_wide_subgroup_intrinsics:1; /* Transform feedback buffer strides in dwords, max. 1K - 4. */ uint8_t xfb_stride[MAX_XFB_BUFFERS]; diff --git a/src/mesa/program/prog_to_nir.c b/src/mesa/program/prog_to_nir.c index 46b73b69b80..3e12b947d00 100644 --- a/src/mesa/program/prog_to_nir.c +++ b/src/mesa/program/prog_to_nir.c @@ -819,6 +819,8 @@ prog_to_nir(const struct gl_context *ctx, const struct gl_program *prog) /* Copy the shader_info from the gl_program */ c->build.shader->info = prog->info; + c->build.shader->info.max_subgroup_size = 128; + c->build.shader->info.min_subgroup_size = 1; s = c->build.shader; diff --git a/src/mesa/state_tracker/st_atifs_to_nir.c b/src/mesa/state_tracker/st_atifs_to_nir.c index 30ff9545997..c534ae008d8 100644 --- a/src/mesa/state_tracker/st_atifs_to_nir.c +++ b/src/mesa/state_tracker/st_atifs_to_nir.c @@ -441,6 +441,8 @@ st_translate_atifs_program(struct ati_fragment_shader *atifs, /* Copy the shader_info from the gl_program */ t->b->shader->info = program->info; + t->b->shader->info.max_subgroup_size = 128; + t->b->shader->info.min_subgroup_size = 1; nir_shader *s = t->b->shader; s->info.name = ralloc_asprintf(s, "ATIFS%d", program->Id);