diff --git a/src/compiler/nir/nir_opt_large_constants.c b/src/compiler/nir/nir_opt_large_constants.c index 963043c08b2..6ef6593b870 100644 --- a/src/compiler/nir/nir_opt_large_constants.c +++ b/src/compiler/nir/nir_opt_large_constants.c @@ -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); } diff --git a/src/compiler/nir/tests/opt_large_constants_tests.cpp b/src/compiler/nir/tests/opt_large_constants_tests.cpp index 47bb8d03c33..f8e010c7330 100644 --- a/src/compiler/nir/tests/opt_large_constants_tests.cpp +++ b/src/compiler/nir/tests/opt_large_constants_tests.cpp @@ -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 {