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

nir: Fix lowering terminology of compute system values: "from"->"to". 

This is to match other NIR terminology.

Signed-off-by: Timur Kristóf <timur.kristof@gmail.com>
Reviewed-by: Jason Ekstrand <jason.ekstrand@collabora.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/15103>
This commit is contained in:
Timur Kristóf 2022-02-24 10:14:08 +01:00 committed by Marge Bot
parent 541f08cd4c
commit 64acec0ef9
11 changed files with 22 additions and 22 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
src/amd/vulkan/radv_shader.c
View File

@ -723,7 +723,7 @@ radv_shader_compile_to_nir(struct radv_device *device, struct vk_shader_module *
/* Mesh shaders run as NGG which can implement local_invocation_index from
* the wave ID in merged_wave_info, but they don't have local_invocation_ids.
*/
.lower_cs_local_id_from_index = nir->info.stage == MESA_SHADER_MESH,
.lower_cs_local_id_to_index = nir->info.stage == MESA_SHADER_MESH,
.lower_local_invocation_index = nir->info.stage == MESA_SHADER_COMPUTE &&
((nir->info.workgroup_size[0] == 1) +
(nir->info.workgroup_size[1] == 1) +

2
src/asahi/compiler/agx_compile.h
View File

@ -216,7 +216,7 @@ static const nir_shader_compiler_options agx_nir_options = {
.lower_insert_byte = true,
.lower_insert_word = true,
.lower_uniforms_to_ubo = true,
.lower_cs_local_index_from_id = true,
.lower_cs_local_index_to_id = true,
.lower_doubles_options = nir_lower_dmod,
.lower_int64_options = ~(nir_lower_iadd64 | nir_lower_imul_2x32_64),

2
src/broadcom/vulkan/v3dv_pipeline.c
View File

@ -204,7 +204,7 @@ const nir_shader_compiler_options v3dv_nir_options = {
.lower_bitfield_extract_to_shifts = true,
.lower_bitfield_reverse = true,
.lower_bit_count = true,
.lower_cs_local_id_from_index = true,
.lower_cs_local_id_to_index = true,
.lower_ffract = true,
.lower_fmod = true,
.lower_pack_unorm_2x16 = true,

6
src/compiler/nir/nir.h
View File

@ -3337,8 +3337,8 @@ typedef struct nir_shader_compiler_options {
*/
bool optimize_sample_mask_in;
bool lower_cs_local_index_from_id;
bool lower_cs_local_id_from_index;
bool lower_cs_local_index_to_id;
bool lower_cs_local_id_to_index;
/* Prevents lowering global_invocation_id to be in terms of workgroup_id */
bool has_cs_global_id;
@ -4763,7 +4763,7 @@ typedef struct nir_lower_compute_system_values_options {
bool has_base_workgroup_id:1;
bool shuffle_local_ids_for_quad_derivatives:1;
bool lower_local_invocation_index:1;
bool lower_cs_local_id_from_index:1;
bool lower_cs_local_id_to_index:1;
} nir_lower_compute_system_values_options;
bool nir_lower_compute_system_values(nir_shader *shader,

16
src/compiler/nir/nir_lower_system_values.c
View File

@ -288,12 +288,12 @@ lower_compute_system_value_instr(nir_builder *b,
switch (intrin->intrinsic) {
case nir_intrinsic_load_local_invocation_id:
/* If lower_cs_local_id_from_index is true, then we derive the local
* index from the local id.
/* If lower_cs_local_id_to_index is true, then we replace
* local_invocation_id with a formula based on local_invocation_index.
*/
if (b->shader->options->lower_cs_local_id_from_index ||
(options && options->lower_cs_local_id_from_index)) {
/* We lower gl_LocalInvocationID from gl_LocalInvocationIndex based
if (b->shader->options->lower_cs_local_id_to_index ||
(options && options->lower_cs_local_id_to_index)) {
/* We lower gl_LocalInvocationID to gl_LocalInvocationIndex based
* on this formula:
*
* gl_LocalInvocationID.x =
@ -422,10 +422,10 @@ lower_compute_system_value_instr(nir_builder *b,
return NULL;
case nir_intrinsic_load_local_invocation_index:
/* If lower_cs_local_index_from_id is true, then we derive the local
* index from the local id.
/* If lower_cs_local_index_to_id is true, then we replace
* local_invocation_index with a formula based on local_invocation_id.
*/
if (b->shader->options->lower_cs_local_index_from_id ||
if (b->shader->options->lower_cs_local_index_to_id ||
(options && options->lower_local_invocation_index)) {
/* From the GLSL man page for gl_LocalInvocationIndex:
*

4
src/freedreno/ir3/ir3_compiler.c
View File

@ -113,7 +113,7 @@ static const nir_shader_compiler_options options = {
.has_fsub = true,
.has_isub = true,
.lower_wpos_pntc = true,
.lower_cs_local_index_from_id = true,
.lower_cs_local_index_to_id = true,
/* Only needed for the spirv_to_nir() pass done in ir3_cmdline.c
* but that should be harmless for GL since 64b is not
@ -172,7 +172,7 @@ static const nir_shader_compiler_options options_a6xx = {
.max_unroll_iterations = 32,
.force_indirect_unrolling = nir_var_all,
.lower_wpos_pntc = true,
.lower_cs_local_index_from_id = true,
.lower_cs_local_index_to_id = true,
/* Only needed for the spirv_to_nir() pass done in ir3_cmdline.c
* but that should be harmless for GL since 64b is not

2
src/gallium/drivers/i915/i915_screen.c
View File

@ -160,7 +160,7 @@ static const struct nir_shader_compiler_options gallivm_nir_options = {
.lower_ifind_msb = true,
.max_unroll_iterations = 32,
.use_interpolated_input_intrinsics = true,
.lower_cs_local_index_from_id = true,
.lower_cs_local_index_to_id = true,
.lower_uniforms_to_ubo = true,
.lower_vector_cmp = true,
.lower_device_index_to_zero = true,

4
src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp
View File

@ -3288,8 +3288,8 @@ nvir_nir_shader_compiler_options(int chipset)
op.lower_base_vertex = false;
op.lower_helper_invocation = false;
op.optimize_sample_mask_in = false;
op.lower_cs_local_index_from_id = true;
op.lower_cs_local_id_from_index = false;
op.lower_cs_local_index_to_id = true;
op.lower_cs_local_id_to_index = false;
op.lower_device_index_to_zero = false; // TODO
op.lower_wpos_pntc = false; // TODO
op.lower_hadd = true; // TODO

2
src/gallium/drivers/v3d/v3d_screen.c
View File

@ -695,7 +695,7 @@ static const nir_shader_compiler_options v3d_nir_options = {
.lower_bitfield_extract_to_shifts = true,
.lower_bitfield_reverse = true,
.lower_bit_count = true,
.lower_cs_local_id_from_index = true,
.lower_cs_local_id_to_index = true,
.lower_ffract = true,
.lower_fmod = true,
.lower_pack_unorm_2x16 = true,

2
src/panfrost/bifrost/bifrost_compile.h
View File

@ -91,7 +91,7 @@ static const nir_shader_compiler_options bifrost_nir_options = {
.has_cs_global_id = true,
.vertex_id_zero_based = true,
.lower_cs_local_index_from_id = true,
.lower_cs_local_index_to_id = true,
.max_unroll_iterations = 32,
.force_indirect_unrolling = (nir_var_shader_in | nir_var_shader_out | nir_var_function_temp),
};

2
src/panfrost/midgard/midgard_compile.h
View File

@ -94,7 +94,7 @@ static const nir_shader_compiler_options midgard_nir_options = {
.vertex_id_zero_based = true,
.has_cs_global_id = true,
.lower_cs_local_index_from_id = true,
.lower_cs_local_index_to_id = true,
.max_unroll_iterations = 32,
.force_indirect_unrolling = (nir_var_shader_in | nir_var_shader_out | nir_var_function_temp),
};