AlexIndustrial/mesa
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

nir: define new subgroup size info

Browse Source 
Reviewed-by: Marek Olšák <marek.olsak@amd.com>
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37258>
This commit is contained in:
Georg Lehmann
2025-09-09 16:41:00 +02:00
committed by Marge Bot
parent 95c2a65662
commit ce91b0be08
9 changed files with 88 additions and 20 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/compiler/nir/nir.c
+2
View File
@@ -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);
src/compiler/nir/nir_print.c
+5
View File
@@ -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];
src/compiler/nir/nir_validate.c
+13
View File
@@ -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);
src/compiler/nir/tests/lower_discard_if_tests.cpp
+7 -4
View File
@@ -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:
}
)"));
}
}
src/compiler/nir/tests/minimize_call_live_states_test.cpp
+10 -5
View File
@@ -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 {
src/compiler/nir/tests/opt_loop_tests.cpp
+18 -9
View File
@@ -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)
src/compiler/shader_info.h
+29 -2
View File
@@ -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];
src/mesa/program/prog_to_nir.c
+2
View File
@@ -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;
src/mesa/state_tracker/st_atifs_to_nir.c
+2
View File
@@ -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);