From 64acec0ef9f640ad27d6fa0fe7622ec63c6900a9 Mon Sep 17 00:00:00 2001 From: =?utf8?q?Timur=20Krist=C3=B3f?= Date: Thu, 24 Feb 2022 10:14:08 +0100 Subject: [PATCH] nir: Fix lowering terminology of compute system values: "from"->"to". MIME-Version: 1.0 Content-Type: text/plain; charset=utf8 Content-Transfer-Encoding: 8bit This is to match other NIR terminology. Signed-off-by: Timur Kristóf Reviewed-by: Jason Ekstrand Reviewed-by: Daniel Schürmann Part-of: --- src/amd/vulkan/radv_shader.c | 2 +- src/asahi/compiler/agx_compile.h | 2 +- src/broadcom/vulkan/v3dv_pipeline.c | 2 +- src/compiler/nir/nir.h | 6 +++--- src/compiler/nir/nir_lower_system_values.c | 16 ++++++++-------- src/freedreno/ir3/ir3_compiler.c | 4 ++-- src/gallium/drivers/i915/i915_screen.c | 2 +- src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp | 4 ++-- src/gallium/drivers/v3d/v3d_screen.c | 2 +- src/panfrost/bifrost/bifrost_compile.h | 2 +- src/panfrost/midgard/midgard_compile.h | 2 +- 11 files changed, 22 insertions(+), 22 deletions(-) diff --git a/src/amd/vulkan/radv_shader.c b/src/amd/vulkan/radv_shader.c index 4a5cae3..b0d3145 100644 --- a/src/amd/vulkan/radv_shader.c +++ b/src/amd/vulkan/radv_shader.c @@ -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) + diff --git a/src/asahi/compiler/agx_compile.h b/src/asahi/compiler/agx_compile.h index f1da074..8df4576 100644 --- a/src/asahi/compiler/agx_compile.h +++ b/src/asahi/compiler/agx_compile.h @@ -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), diff --git a/src/broadcom/vulkan/v3dv_pipeline.c b/src/broadcom/vulkan/v3dv_pipeline.c index 0f4d43d..31d5b6e 100644 --- a/src/broadcom/vulkan/v3dv_pipeline.c +++ b/src/broadcom/vulkan/v3dv_pipeline.c @@ -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, diff --git a/src/compiler/nir/nir.h b/src/compiler/nir/nir.h index a48fbb2..a1904e7 100644 --- a/src/compiler/nir/nir.h +++ b/src/compiler/nir/nir.h @@ -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, diff --git a/src/compiler/nir/nir_lower_system_values.c b/src/compiler/nir/nir_lower_system_values.c index 37657b0..a2da7e5 100644 --- a/src/compiler/nir/nir_lower_system_values.c +++ b/src/compiler/nir/nir_lower_system_values.c @@ -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: * diff --git a/src/freedreno/ir3/ir3_compiler.c b/src/freedreno/ir3/ir3_compiler.c index c23d85e..3c9e0db 100644 --- a/src/freedreno/ir3/ir3_compiler.c +++ b/src/freedreno/ir3/ir3_compiler.c @@ -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 diff --git a/src/gallium/drivers/i915/i915_screen.c b/src/gallium/drivers/i915/i915_screen.c index 5ef24df..cfe1096 100644 --- a/src/gallium/drivers/i915/i915_screen.c +++ b/src/gallium/drivers/i915/i915_screen.c @@ -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, diff --git a/src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp b/src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp index 45cf808..87f6b67 100644 --- a/src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp +++ b/src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp @@ -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 diff --git a/src/gallium/drivers/v3d/v3d_screen.c b/src/gallium/drivers/v3d/v3d_screen.c index 6aa4568..955e198 100644 --- a/src/gallium/drivers/v3d/v3d_screen.c +++ b/src/gallium/drivers/v3d/v3d_screen.c @@ -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, diff --git a/src/panfrost/bifrost/bifrost_compile.h b/src/panfrost/bifrost/bifrost_compile.h index 9b320d7..446ef24 100644 --- a/src/panfrost/bifrost/bifrost_compile.h +++ b/src/panfrost/bifrost/bifrost_compile.h @@ -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), }; diff --git a/src/panfrost/midgard/midgard_compile.h b/src/panfrost/midgard/midgard_compile.h index b24c63b..f9c44f9 100644 --- a/src/panfrost/midgard/midgard_compile.h +++ b/src/panfrost/midgard/midgard_compile.h @@ -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), }; -- 2.7.4