dsg/mesa
2
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

nir/opt_large_constants: don't add constants implemented with ALU to the constant data

Browse source 
Reviewed-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/33002>
This commit is contained in:
Georg Lehmann 2026-03-18 11:58:46 +01:00 • committed by Marge Bot
parent 581df90a89
commit e810382a1e
2 changed files with 9 additions and 12 deletions
Download patch file Download diff file Expand all files Collapse all files

16
src/compiler/nir/nir_opt_large_constants.c
View file

@ -498,12 +498,14 @@ nir_opt_large_constants(nir_shader *shader,
continue;
}
if (i > 0 && var_info_cmp(info, &var_infos[i - 1]) == 0) {
info->var->data.location = var_infos[i - 1].var->data.location;
info->duplicate = true;
} else {
info->var->data.location = ALIGN_POT(shader->constant_data_size, var_align);
shader->constant_data_size = info->var->data.location + var_size;
if (!info->is_small) {
if (i > 0 && var_info_cmp(info, &var_infos[i - 1]) == 0) {
info->var->data.location = var_infos[i - 1].var->data.location;
info->duplicate = true;
} else {
info->var->data.location = ALIGN_POT(shader->constant_data_size, var_align);
shader->constant_data_size = info->var->data.location + var_size;
}
}
has_constant |= info->is_constant;
@ -522,7 +524,7 @@ nir_opt_large_constants(nir_shader *shader,
shader->constant_data_size);
for (int i = 0; i < num_locals; i++) {
struct var_info *info = &var_infos[i];
if (!info->duplicate && info->is_constant) {
if (!info->duplicate && info->is_constant && !info->is_small) {
memcpy((char *)shader->constant_data + info->var->data.location,
info->constant_data, info->constant_data_size);
}

5
src/compiler/nir/tests/opt_large_constants_tests.cpp
View file

@ -63,7 +63,6 @@ TEST_F(nir_large_constants_test, small_int_array)
workgroup_size: 1, 1, 1
max_subgroup_size: 128
min_subgroup_size: 1
constants: 32
decl_function main () (entrypoint)
impl main {
@ -97,7 +96,6 @@ TEST_F(nir_large_constants_test, small_uint8_t_array)
workgroup_size: 1, 1, 1
max_subgroup_size: 128
min_subgroup_size: 1
constants: 8
decl_function main () (entrypoint)
impl main {
@ -132,7 +130,6 @@ TEST_F(nir_large_constants_test, small_bool_array)
workgroup_size: 1, 1, 1
max_subgroup_size: 128
min_subgroup_size: 1
constants: 32
decl_function main () (entrypoint)
impl main {
@ -166,7 +163,6 @@ TEST_F(nir_large_constants_test, small_uint64_t_array)
workgroup_size: 1, 1, 1
max_subgroup_size: 128
min_subgroup_size: 1
constants: 64
decl_function main () (entrypoint)
impl main {
@ -201,7 +197,6 @@ TEST_F(nir_large_constants_test, small_float_natural_numbers_including_zero_arra
workgroup_size: 1, 1, 1
max_subgroup_size: 128
min_subgroup_size: 1
constants: 32
decl_function main () (entrypoint)
impl main {