1/*
2 * Copyright © 2017 Ilia Mirkin
3 *
4 * Permission is hereby granted, free of charge, to any person obtaining a
5 * copy of this software and associated documentation files (the "Software"),
6 * to deal in the Software without restriction, including without limitation
7 * the rights to use, copy, modify, merge, publish, distribute, sublicense,
8 * and/or sell copies of the Software, and to permit persons to whom the
9 * Software is furnished to do so, subject to the following conditions:
10 *
11 * The above copyright notice and this permission notice (including the next
12 * paragraph) shall be included in all copies or substantial portions of the
13 * Software.
14 *
15 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
16 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
17 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL
18 * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
19 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
20 * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER
21 * DEALINGS IN THE SOFTWARE.
22 */
23
24/**
25 * \file lower_cs_derived.cpp
26 *
27 * For hardware that does not support the gl_GlobalInvocationID and
28 * gl_LocalInvocationIndex system values, replace them with fresh
29 * globals. Note that we can't rely on gl_WorkGroupSize or
30 * gl_LocalGroupSizeARB being available, since they may only have been defined
31 * in a non-main shader.
32 *
33 * [ This can happen if only a secondary shader has the layout(local_size_*)
34 *   declaration. ]
35 *
36 * This is meant to be run post-linking.
37 */
38
39#include "glsl_symbol_table.h"
40#include "ir_hierarchical_visitor.h"
41#include "ir.h"
42#include "ir_builder.h"
43#include "linker.h"
44#include "program/prog_statevars.h"
45#include "builtin_functions.h"
46#include "main/shader_types.h"
47
48using namespace ir_builder;
49
50namespace {
51
52class lower_cs_derived_visitor : public ir_hierarchical_visitor {
53public:
54   explicit lower_cs_derived_visitor(gl_linked_shader *shader)
55      : progress(false),
56        shader(shader),
57        local_size_variable(shader->Program->info.workgroup_size_variable),
58        gl_WorkGroupSize(NULL),
59        gl_WorkGroupID(NULL),
60        gl_LocalInvocationID(NULL),
61        gl_GlobalInvocationID(NULL),
62        gl_LocalInvocationIndex(NULL)
63   {
64      main_sig = _mesa_get_main_function_signature(shader->symbols);
65      assert(main_sig);
66   }
67
68   virtual ir_visitor_status visit(ir_dereference_variable *);
69
70   ir_variable *add_system_value(
71         int slot, const glsl_type *type, const char *name);
72   void find_sysvals();
73   void make_gl_GlobalInvocationID();
74   void make_gl_LocalInvocationIndex();
75
76   bool progress;
77
78private:
79   gl_linked_shader *shader;
80   bool local_size_variable;
81   ir_function_signature *main_sig;
82
83   ir_rvalue *gl_WorkGroupSize;
84   ir_variable *gl_WorkGroupID;
85   ir_variable *gl_LocalInvocationID;
86
87   ir_variable *gl_GlobalInvocationID;
88   ir_variable *gl_LocalInvocationIndex;
89};
90
91} /* anonymous namespace */
92
93ir_variable *
94lower_cs_derived_visitor::add_system_value(
95      int slot, const glsl_type *type, const char *name)
96{
97   ir_variable *var = new(shader) ir_variable(type, name, ir_var_system_value);
98   var->data.how_declared = ir_var_declared_implicitly;
99   var->data.read_only = true;
100   var->data.location = slot;
101   var->data.explicit_location = true;
102   var->data.explicit_index = 0;
103   shader->ir->push_head(var);
104
105   return var;
106}
107
108void
109lower_cs_derived_visitor::find_sysvals()
110{
111   if (gl_WorkGroupSize != NULL)
112      return;
113
114   ir_variable *WorkGroupSize;
115   if (local_size_variable)
116      WorkGroupSize = shader->symbols->get_variable("gl_LocalGroupSizeARB");
117   else
118      WorkGroupSize = shader->symbols->get_variable("gl_WorkGroupSize");
119   if (WorkGroupSize)
120      gl_WorkGroupSize = new(shader) ir_dereference_variable(WorkGroupSize);
121   gl_WorkGroupID = shader->symbols->get_variable("gl_WorkGroupID");
122   gl_LocalInvocationID = shader->symbols->get_variable("gl_LocalInvocationID");
123
124   /*
125    * These may be missing due to either dead code elimination, or, in the
126    * case of the group size, due to the layout being declared in a non-main
127    * shader. Re-create them.
128    */
129
130   if (!gl_WorkGroupID)
131      gl_WorkGroupID = add_system_value(
132            SYSTEM_VALUE_WORKGROUP_ID, glsl_type::uvec3_type, "gl_WorkGroupID");
133   if (!gl_LocalInvocationID)
134      gl_LocalInvocationID = add_system_value(
135            SYSTEM_VALUE_LOCAL_INVOCATION_ID, glsl_type::uvec3_type,
136            "gl_LocalInvocationID");
137   if (!WorkGroupSize) {
138      if (local_size_variable) {
139         gl_WorkGroupSize = new(shader) ir_dereference_variable(
140               add_system_value(
141                     SYSTEM_VALUE_WORKGROUP_SIZE, glsl_type::uvec3_type,
142                     "gl_LocalGroupSizeARB"));
143      } else {
144         ir_constant_data data;
145         memset(&data, 0, sizeof(data));
146         for (int i = 0; i < 3; i++)
147            data.u[i] = shader->Program->info.workgroup_size[i];
148         gl_WorkGroupSize = new(shader) ir_constant(glsl_type::uvec3_type, &data);
149      }
150   }
151}
152
153void
154lower_cs_derived_visitor::make_gl_GlobalInvocationID()
155{
156   if (gl_GlobalInvocationID != NULL)
157      return;
158
159   find_sysvals();
160
161   /* gl_GlobalInvocationID =
162    *    gl_WorkGroupID * gl_WorkGroupSize + gl_LocalInvocationID
163    */
164   gl_GlobalInvocationID = new(shader) ir_variable(
165         glsl_type::uvec3_type, "__GlobalInvocationID", ir_var_temporary);
166   shader->ir->push_head(gl_GlobalInvocationID);
167
168   ir_instruction *inst =
169      assign(gl_GlobalInvocationID,
170             add(mul(gl_WorkGroupID, gl_WorkGroupSize->clone(shader, NULL)),
171                 gl_LocalInvocationID));
172   main_sig->body.push_head(inst);
173}
174
175void
176lower_cs_derived_visitor::make_gl_LocalInvocationIndex()
177{
178   if (gl_LocalInvocationIndex != NULL)
179      return;
180
181   find_sysvals();
182
183   /* gl_LocalInvocationIndex =
184    *    gl_LocalInvocationID.z * gl_WorkGroupSize.x * gl_WorkGroupSize.y +
185    *    gl_LocalInvocationID.y * gl_WorkGroupSize.x +
186    *    gl_LocalInvocationID.x;
187    */
188   gl_LocalInvocationIndex = new(shader)
189      ir_variable(glsl_type::uint_type, "__LocalInvocationIndex", ir_var_temporary);
190   shader->ir->push_head(gl_LocalInvocationIndex);
191
192   ir_expression *index_z =
193      mul(mul(swizzle_z(gl_LocalInvocationID), swizzle_x(gl_WorkGroupSize->clone(shader, NULL))),
194          swizzle_y(gl_WorkGroupSize->clone(shader, NULL)));
195   ir_expression *index_y =
196      mul(swizzle_y(gl_LocalInvocationID), swizzle_x(gl_WorkGroupSize->clone(shader, NULL)));
197   ir_expression *index_y_plus_z = add(index_y, index_z);
198   operand index_x(swizzle_x(gl_LocalInvocationID));
199   ir_expression *index_x_plus_y_plus_z = add(index_y_plus_z, index_x);
200   ir_instruction *inst =
201      assign(gl_LocalInvocationIndex, index_x_plus_y_plus_z);
202   main_sig->body.push_head(inst);
203}
204
205ir_visitor_status
206lower_cs_derived_visitor::visit(ir_dereference_variable *ir)
207{
208   if (ir->var->data.mode == ir_var_system_value &&
209       ir->var->data.location == SYSTEM_VALUE_GLOBAL_INVOCATION_ID) {
210      make_gl_GlobalInvocationID();
211      ir->var = gl_GlobalInvocationID;
212      progress = true;
213   }
214
215   if (ir->var->data.mode == ir_var_system_value &&
216       ir->var->data.location == SYSTEM_VALUE_LOCAL_INVOCATION_INDEX) {
217      make_gl_LocalInvocationIndex();
218      ir->var = gl_LocalInvocationIndex;
219      progress = true;
220   }
221
222   return visit_continue;
223}
224
225bool
226lower_cs_derived(gl_linked_shader *shader)
227{
228   if (shader->Stage != MESA_SHADER_COMPUTE)
229      return false;
230
231   lower_cs_derived_visitor v(shader);
232   v.run(shader->ir);
233
234   return v.progress;
235}
236