mesa: Remove PIPE_CAP_CS_DERIVED_SYSTEM_VALUES_SUPPORTED/lower_cs_derived.

We have fine NIR lowering for this (already called from mesa/st), no need
for a separate GLSL pass.

Reviewed-by: Timothy Arceri <tarceri@itsqueeze.com>
Reviewed-by: Alyssa Rosenzweig <alyssa.rosenzweig@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/18361>

docs/gallium/screen.rst

@@ -561,9 +561,6 @@ The integer capabilities:
   OpenMAX should use a compute-based blit instead of pipe_context::blit and compute pipeline for compositing images.
 * ``PIPE_CAP_FRAGMENT_SHADER_INTERLOCK``: True if fragment shader interlock
   functionality is supported.
-* ``PIPE_CAP_CS_DERIVED_SYSTEM_VALUES_SUPPORTED``: True if driver handles
-  gl_LocalInvocationIndex and gl_GlobalInvocationID.  Otherwise, gallium frontends will
-  lower those system values.
 * ``PIPE_CAP_ATOMIC_FLOAT_MINMAX``: Atomic float point minimum,
   maximum, exchange and compare-and-swap support to buffer and shared variables.
 * ``PIPE_CAP_TGSI_DIV``: Whether opcode DIV is supported

src/compiler/glsl/ir_optimization.h

@@ -122,7 +122,6 @@ void optimize_dead_builtin_variables(exec_list *instructions,
                                      enum ir_variable_mode other);
 bool lower_tess_level(gl_linked_shader *shader);
 
-bool lower_cs_derived(gl_linked_shader *shader);
 bool lower_blend_equation_advanced(gl_linked_shader *shader, bool coherent);
 
 bool lower_builtins(exec_list *instructions);

src/compiler/glsl/linker.cpp

@@ -2544,9 +2544,6 @@ link_intrastage_shaders(void *mem_ctx,
       }
    }
 
-   if (ctx->Const.LowerCsDerivedVariables)
-      lower_cs_derived(linked);
-
    /* Set the linked source SHA1. */
    if (num_shaders == 1) {
       memcpy(linked->linked_source_sha1, shader_list[0]->compiled_source_sha1,

src/compiler/glsl/lower_cs_derived.cpp

@@ -1,235 +0,0 @@
-/*
- * Copyright © 2017 Ilia Mirkin
- *
- * Permission is hereby granted, free of charge, to any person obtaining a
- * copy of this software and associated documentation files (the "Software"),
- * to deal in the Software without restriction, including without limitation
- * the rights to use, copy, modify, merge, publish, distribute, sublicense,
- * and/or sell copies of the Software, and to permit persons to whom the
- * Software is furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice (including the next
- * paragraph) shall be included in all copies or substantial portions of the
- * Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL
- * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
- * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER
- * DEALINGS IN THE SOFTWARE.
- */
-
-/**
- * \file lower_cs_derived.cpp
- *
- * For hardware that does not support the gl_GlobalInvocationID and
- * gl_LocalInvocationIndex system values, replace them with fresh
- * globals. Note that we can't rely on gl_WorkGroupSize or
- * gl_LocalGroupSizeARB being available, since they may only have been defined
- * in a non-main shader.
- *
- * [ This can happen if only a secondary shader has the layout(local_size_*)
- *   declaration. ]
- *
- * This is meant to be run post-linking.
- */
-
-#include "glsl_symbol_table.h"
-#include "ir_hierarchical_visitor.h"
-#include "ir.h"
-#include "ir_builder.h"
-#include "linker.h"
-#include "program/prog_statevars.h"
-#include "builtin_functions.h"
-#include "main/shader_types.h"
-
-using namespace ir_builder;
-
-namespace {
-
-class lower_cs_derived_visitor : public ir_hierarchical_visitor {
-public:
-   explicit lower_cs_derived_visitor(gl_linked_shader *shader)
-      : progress(false),
-        shader(shader),
-        local_size_variable(shader->Program->info.workgroup_size_variable),
-        gl_WorkGroupSize(NULL),
-        gl_WorkGroupID(NULL),
-        gl_LocalInvocationID(NULL),
-        gl_GlobalInvocationID(NULL),
-        gl_LocalInvocationIndex(NULL)
-   {
-      main_sig = _mesa_get_main_function_signature(shader->symbols);
-      assert(main_sig);
-   }
-
-   virtual ir_visitor_status visit(ir_dereference_variable *);
-
-   ir_variable *add_system_value(
-         int slot, const glsl_type *type, const char *name);
-   void find_sysvals();
-   void make_gl_GlobalInvocationID();
-   void make_gl_LocalInvocationIndex();
-
-   bool progress;
-
-private:
-   gl_linked_shader *shader;
-   bool local_size_variable;
-   ir_function_signature *main_sig;
-
-   ir_rvalue *gl_WorkGroupSize;
-   ir_variable *gl_WorkGroupID;
-   ir_variable *gl_LocalInvocationID;
-
-   ir_variable *gl_GlobalInvocationID;
-   ir_variable *gl_LocalInvocationIndex;
-};
-
-} /* anonymous namespace */
-
-ir_variable *
-lower_cs_derived_visitor::add_system_value(
-      int slot, const glsl_type *type, const char *name)
-{
-   ir_variable *var = new(shader) ir_variable(type, name, ir_var_system_value);
-   var->data.how_declared = ir_var_declared_implicitly;
-   var->data.read_only = true;
-   var->data.location = slot;
-   var->data.explicit_location = true;
-   var->data.explicit_index = 0;
-   shader->ir->push_head(var);
-
-   return var;
-}
-
-void
-lower_cs_derived_visitor::find_sysvals()
-{
-   if (gl_WorkGroupSize != NULL)
-      return;
-
-   ir_variable *WorkGroupSize;
-   if (local_size_variable)
-      WorkGroupSize = shader->symbols->get_variable("gl_LocalGroupSizeARB");
-   else
-      WorkGroupSize = shader->symbols->get_variable("gl_WorkGroupSize");
-   if (WorkGroupSize)
-      gl_WorkGroupSize = new(shader) ir_dereference_variable(WorkGroupSize);
-   gl_WorkGroupID = shader->symbols->get_variable("gl_WorkGroupID");
-   gl_LocalInvocationID = shader->symbols->get_variable("gl_LocalInvocationID");
-
-   /*
-    * These may be missing due to either dead code elimination, or, in the
-    * case of the group size, due to the layout being declared in a non-main
-    * shader. Re-create them.
-    */
-
-   if (!gl_WorkGroupID)
-      gl_WorkGroupID = add_system_value(
-            SYSTEM_VALUE_WORKGROUP_ID, glsl_type::uvec3_type, "gl_WorkGroupID");
-   if (!gl_LocalInvocationID)
-      gl_LocalInvocationID = add_system_value(
-            SYSTEM_VALUE_LOCAL_INVOCATION_ID, glsl_type::uvec3_type,
-            "gl_LocalInvocationID");
-   if (!WorkGroupSize) {
-      if (local_size_variable) {
-         gl_WorkGroupSize = new(shader) ir_dereference_variable(
-               add_system_value(
-                     SYSTEM_VALUE_WORKGROUP_SIZE, glsl_type::uvec3_type,
-                     "gl_LocalGroupSizeARB"));
-      } else {
-         ir_constant_data data;
-         memset(&data, 0, sizeof(data));
-         for (int i = 0; i < 3; i++)
-            data.u[i] = shader->Program->info.workgroup_size[i];
-         gl_WorkGroupSize = new(shader) ir_constant(glsl_type::uvec3_type, &data);
-      }
-   }
-}
-
-void
-lower_cs_derived_visitor::make_gl_GlobalInvocationID()
-{
-   if (gl_GlobalInvocationID != NULL)
-      return;
-
-   find_sysvals();
-
-   /* gl_GlobalInvocationID =
-    *    gl_WorkGroupID * gl_WorkGroupSize + gl_LocalInvocationID
-    */
-   gl_GlobalInvocationID = new(shader) ir_variable(
-         glsl_type::uvec3_type, "__GlobalInvocationID", ir_var_temporary);
-   shader->ir->push_head(gl_GlobalInvocationID);
-
-   ir_instruction *inst =
-      assign(gl_GlobalInvocationID,
-             add(mul(gl_WorkGroupID, gl_WorkGroupSize->clone(shader, NULL)),
-                 gl_LocalInvocationID));
-   main_sig->body.push_head(inst);
-}
-
-void
-lower_cs_derived_visitor::make_gl_LocalInvocationIndex()
-{
-   if (gl_LocalInvocationIndex != NULL)
-      return;
-
-   find_sysvals();
-
-   /* gl_LocalInvocationIndex =
-    *    gl_LocalInvocationID.z * gl_WorkGroupSize.x * gl_WorkGroupSize.y +
-    *    gl_LocalInvocationID.y * gl_WorkGroupSize.x +
-    *    gl_LocalInvocationID.x;
-    */
-   gl_LocalInvocationIndex = new(shader)
-      ir_variable(glsl_type::uint_type, "__LocalInvocationIndex", ir_var_temporary);
-   shader->ir->push_head(gl_LocalInvocationIndex);
-
-   ir_expression *index_z =
-      mul(mul(swizzle_z(gl_LocalInvocationID), swizzle_x(gl_WorkGroupSize->clone(shader, NULL))),
-          swizzle_y(gl_WorkGroupSize->clone(shader, NULL)));
-   ir_expression *index_y =
-      mul(swizzle_y(gl_LocalInvocationID), swizzle_x(gl_WorkGroupSize->clone(shader, NULL)));
-   ir_expression *index_y_plus_z = add(index_y, index_z);
-   operand index_x(swizzle_x(gl_LocalInvocationID));
-   ir_expression *index_x_plus_y_plus_z = add(index_y_plus_z, index_x);
-   ir_instruction *inst =
-      assign(gl_LocalInvocationIndex, index_x_plus_y_plus_z);
-   main_sig->body.push_head(inst);
-}
-
-ir_visitor_status
-lower_cs_derived_visitor::visit(ir_dereference_variable *ir)
-{
-   if (ir->var->data.mode == ir_var_system_value &&
-       ir->var->data.location == SYSTEM_VALUE_GLOBAL_INVOCATION_ID) {
-      make_gl_GlobalInvocationID();
-      ir->var = gl_GlobalInvocationID;
-      progress = true;
-   }
-
-   if (ir->var->data.mode == ir_var_system_value &&
-       ir->var->data.location == SYSTEM_VALUE_LOCAL_INVOCATION_INDEX) {
-      make_gl_LocalInvocationIndex();
-      ir->var = gl_LocalInvocationIndex;
-      progress = true;
-   }
-
-   return visit_continue;
-}
-
-bool
-lower_cs_derived(gl_linked_shader *shader)
-{
-   if (shader->Stage != MESA_SHADER_COMPUTE)
-      return false;
-
-   lower_cs_derived_visitor v(shader);
-   v.run(shader->ir);
-
-   return v.progress;
-}

src/compiler/glsl/meson.build

@@ -156,7 +156,6 @@ files_libglsl = files(
   'list.h',
   'lower_blend_equation_advanced.cpp',
   'lower_builtins.cpp',
-  'lower_cs_derived.cpp',
   'lower_discard.cpp',
   'lower_discard_flow.cpp',
   'lower_distance.cpp',

src/gallium/auxiliary/nir/nir_to_tgsi.c

@@ -4035,6 +4035,11 @@ static const nir_shader_compiler_options nir_to_tgsi_compiler_options = {
    .lower_vector_cmp = true,
    .lower_int64_options = nir_lower_imul_2x32_64,
    .use_interpolated_input_intrinsics = true,
+
+   /* TGSI doesn't have a semantic for local or global index, just local and
+    * workgroup id.
+    */
+   .lower_cs_local_index_to_id = true,
 };
 
 /* Returns a default compiler options for drivers with only nir-to-tgsi-based

src/gallium/auxiliary/util/u_screen.c

@@ -370,7 +370,6 @@ u_pipe_screen_get_param_defaults(struct pipe_screen *pscreen,
    case PIPE_CAP_IMAGE_STORE_FORMATTED:
    case PIPE_CAP_PREFER_COMPUTE_FOR_MULTIMEDIA:
    case PIPE_CAP_FRAGMENT_SHADER_INTERLOCK:
-   case PIPE_CAP_CS_DERIVED_SYSTEM_VALUES_SUPPORTED:
    case PIPE_CAP_ATOMIC_FLOAT_MINMAX:
    case PIPE_CAP_SHADER_SAMPLES_IDENTICAL:
    case PIPE_CAP_IMAGE_ATOMIC_INC_WRAP:

src/gallium/drivers/asahi/agx_pipe.c

@@ -756,7 +756,6 @@ agx_get_param(struct pipe_screen* pscreen, enum pipe_cap param)
    case PIPE_CAP_TEXTURE_FLOAT_LINEAR:
    case PIPE_CAP_TEXTURE_HALF_FLOAT_LINEAR:
    case PIPE_CAP_SHADER_ARRAY_COMPONENTS:
-   case PIPE_CAP_CS_DERIVED_SYSTEM_VALUES_SUPPORTED:
    case PIPE_CAP_PACKED_UNIFORMS:
       return 1;
 

src/gallium/drivers/crocus/crocus_screen.c

@@ -192,7 +192,6 @@ crocus_get_param(struct pipe_screen *pscreen, enum pipe_cap param)
    case PIPE_CAP_FS_FACE_IS_INTEGER_SYSVAL:
    case PIPE_CAP_INVALIDATE_BUFFER:
    case PIPE_CAP_SURFACE_REINTERPRET_BLOCKS:
-   case PIPE_CAP_CS_DERIVED_SYSTEM_VALUES_SUPPORTED:
    case PIPE_CAP_FENCE_SIGNAL:
    case PIPE_CAP_DEMOTE_TO_HELPER_INVOCATION:
    case PIPE_CAP_GL_CLAMP:

src/gallium/drivers/iris/iris_screen.c

@@ -277,7 +277,6 @@ iris_get_param(struct pipe_screen *pscreen, enum pipe_cap param)
    case PIPE_CAP_COMPUTE_SHADER_DERIVATIVES:
    case PIPE_CAP_INVALIDATE_BUFFER:
    case PIPE_CAP_SURFACE_REINTERPRET_BLOCKS:
-   case PIPE_CAP_CS_DERIVED_SYSTEM_VALUES_SUPPORTED:
    case PIPE_CAP_TEXTURE_SHADOW_LOD:
    case PIPE_CAP_SHADER_SAMPLES_IDENTICAL:
    case PIPE_CAP_GL_SPIRV:

src/gallium/drivers/nouveau/nv50/nv50_screen.c

@@ -381,7 +381,6 @@ nv50_screen_get_param(struct pipe_screen *pscreen, enum pipe_cap param)
    case PIPE_CAP_ATOMIC_FLOAT_MINMAX:
    case PIPE_CAP_CONSERVATIVE_RASTER_INNER_COVERAGE:
    case PIPE_CAP_FRAGMENT_SHADER_INTERLOCK:
-   case PIPE_CAP_CS_DERIVED_SYSTEM_VALUES_SUPPORTED:
    case PIPE_CAP_FBFETCH_COHERENT:
    case PIPE_CAP_IMAGE_ATOMIC_INC_WRAP:
    case PIPE_CAP_DEMOTE_TO_HELPER_INVOCATION:

src/gallium/drivers/nouveau/nvc0/nvc0_screen.c

@@ -409,7 +409,6 @@ nvc0_screen_get_param(struct pipe_screen *pscreen, enum pipe_cap param)
    case PIPE_CAP_ATOMIC_FLOAT_MINMAX:
    case PIPE_CAP_CONSERVATIVE_RASTER_INNER_COVERAGE:
    case PIPE_CAP_FRAGMENT_SHADER_INTERLOCK:
-   case PIPE_CAP_CS_DERIVED_SYSTEM_VALUES_SUPPORTED:
    case PIPE_CAP_FBFETCH_COHERENT:
    case PIPE_CAP_TGSI_TG4_COMPONENT_IN_SWIZZLE:
    case PIPE_CAP_OPENCL_INTEGER_FUNCTIONS: /* could be done */

src/gallium/drivers/panfrost/pan_screen.c

@@ -166,7 +166,6 @@ panfrost_get_param(struct pipe_screen *screen, enum pipe_cap param)
         case PIPE_CAP_TEXTURE_FLOAT_LINEAR:
         case PIPE_CAP_TEXTURE_HALF_FLOAT_LINEAR:
         case PIPE_CAP_SHADER_ARRAY_COMPONENTS:
-        case PIPE_CAP_CS_DERIVED_SYSTEM_VALUES_SUPPORTED:
         case PIPE_CAP_TEXTURE_BUFFER_OBJECTS:
         case PIPE_CAP_TEXTURE_BUFFER_SAMPLER:
         case PIPE_CAP_PACKED_UNIFORMS:

src/gallium/drivers/softpipe/sp_screen.c

@@ -91,6 +91,11 @@ static const nir_shader_compiler_options sp_compiler_options = {
    .lower_int64_options = nir_lower_imul_2x32_64,
    .max_unroll_iterations = 32,
    .use_interpolated_input_intrinsics = true,
+
+   /* TGSI doesn't have a semantic for local or global index, just local and
+    * workgroup id.
+    */
+   .lower_cs_local_index_to_id = true,
 };
 
 static const void *

src/gallium/drivers/svga/svga_screen.c

@@ -761,6 +761,7 @@ vgpu10_get_shader_param(struct pipe_screen *screen,
    .lower_rotate = true,                                                      \
    .lower_uniforms_to_ubo = true,                                             \
    .lower_vector_cmp = true,                                                  \
+   .lower_cs_local_index_to_id = true,                                        \
    .max_unroll_iterations = 32,                                               \
    .use_interpolated_input_intrinsics = true
 

src/gallium/include/pipe/p_defines.h

@@ -938,7 +938,6 @@ enum pipe_cap
    PIPE_CAP_PREFER_COMPUTE_FOR_MULTIMEDIA,
    PIPE_CAP_FRAGMENT_SHADER_INTERLOCK,
    PIPE_CAP_FBFETCH_COHERENT,
-   PIPE_CAP_CS_DERIVED_SYSTEM_VALUES_SUPPORTED,
    PIPE_CAP_ATOMIC_FLOAT_MINMAX,
    PIPE_CAP_TGSI_DIV,
    PIPE_CAP_FRAGMENT_SHADER_TEXTURE_LOD,

src/mesa/main/consts_exts.h

@@ -907,9 +907,6 @@ struct gl_constants
    GLuint MaxTessControlTotalOutputComponents;
    bool LowerTessLevel; /**< Lower gl_TessLevel* from float[n] to vecn? */
    bool PrimitiveRestartForPatches;
-   bool LowerCsDerivedVariables;    /**< Lower gl_GlobalInvocationID and
-                                     *   gl_LocalInvocationIndex based on
-                                     *   other builtin variables. */
 
    /** GL_OES_primitive_bounding_box */
    bool NoPrimitiveBoundingBoxOutput;

src/mesa/state_tracker/st_extensions.c

@@ -392,8 +392,6 @@ void st_init_limits(struct pipe_screen *screen,
       screen->get_param(screen, PIPE_CAP_GLSL_TESS_LEVELS_AS_INPUTS);
    c->LowerTessLevel =
       !screen->get_param(screen, PIPE_CAP_NIR_COMPACT_ARRAYS);
-   c->LowerCsDerivedVariables =
-      !screen->get_param(screen, PIPE_CAP_CS_DERIVED_SYSTEM_VALUES_SUPPORTED);
    c->PrimitiveRestartForPatches = false;
 
    c->MaxCombinedTextureImageUnits =