1 /*
2 * Copyright © 2014 Intel Corporation
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 #include "util/glheader.h"
25 #include "bufferobj.h"
26 #include "context.h"
27 #include "state.h"
28 #include "api_exec_decl.h"
29
30 #include "pipe/p_state.h"
31
32 #include "state_tracker/st_context.h"
33 #include "state_tracker/st_cb_bitmap.h"
34 #include "state_tracker/st_util.h"
35
36 static bool
check_valid_to_compute(struct gl_context * ctx,const char * function)37 check_valid_to_compute(struct gl_context *ctx, const char *function)
38 {
39 if (!_mesa_has_compute_shaders(ctx)) {
40 _mesa_error(ctx, GL_INVALID_OPERATION,
41 "unsupported function (%s) called",
42 function);
43 return false;
44 }
45
46 /* From the OpenGL 4.3 Core Specification, Chapter 19, Compute Shaders:
47 *
48 * "An INVALID_OPERATION error is generated if there is no active program
49 * for the compute shader stage."
50 */
51 if (ctx->_Shader->CurrentProgram[MESA_SHADER_COMPUTE] == NULL) {
52 _mesa_error(ctx, GL_INVALID_OPERATION,
53 "%s(no active compute shader)",
54 function);
55 return false;
56 }
57
58 return true;
59 }
60
61 static bool
validate_DispatchCompute(struct gl_context * ctx,struct pipe_grid_info * info)62 validate_DispatchCompute(struct gl_context *ctx, struct pipe_grid_info *info)
63 {
64 if (!check_valid_to_compute(ctx, "glDispatchCompute"))
65 return GL_FALSE;
66
67 for (int i = 0; i < 3; i++) {
68 /* From the OpenGL 4.3 Core Specification, Chapter 19, Compute Shaders:
69 *
70 * "An INVALID_VALUE error is generated if any of num_groups_x,
71 * num_groups_y and num_groups_z are greater than or equal to the
72 * maximum work group count for the corresponding dimension."
73 *
74 * However, the "or equal to" portions appears to be a specification
75 * bug. In all other areas, the specification appears to indicate that
76 * the number of workgroups can match the MAX_COMPUTE_WORK_GROUP_COUNT
77 * value. For example, under DispatchComputeIndirect:
78 *
79 * "If any of num_groups_x, num_groups_y or num_groups_z is greater than
80 * the value of MAX_COMPUTE_WORK_GROUP_COUNT for the corresponding
81 * dimension then the results are undefined."
82 *
83 * Additionally, the OpenGLES 3.1 specification does not contain "or
84 * equal to" as an error condition.
85 */
86 if (info->grid[i] > ctx->Const.MaxComputeWorkGroupCount[i]) {
87 _mesa_error(ctx, GL_INVALID_VALUE,
88 "glDispatchCompute(num_groups_%c)", 'x' + i);
89 return GL_FALSE;
90 }
91 }
92
93 /* The ARB_compute_variable_group_size spec says:
94 *
95 * "An INVALID_OPERATION error is generated by DispatchCompute if the active
96 * program for the compute shader stage has a variable work group size."
97 */
98 struct gl_program *prog = ctx->_Shader->CurrentProgram[MESA_SHADER_COMPUTE];
99 if (prog->info.workgroup_size_variable) {
100 _mesa_error(ctx, GL_INVALID_OPERATION,
101 "glDispatchCompute(variable work group size forbidden)");
102 return GL_FALSE;
103 }
104
105 return GL_TRUE;
106 }
107
108 static bool
validate_DispatchComputeGroupSizeARB(struct gl_context * ctx,struct pipe_grid_info * info)109 validate_DispatchComputeGroupSizeARB(struct gl_context *ctx,
110 struct pipe_grid_info *info)
111 {
112 if (!check_valid_to_compute(ctx, "glDispatchComputeGroupSizeARB"))
113 return GL_FALSE;
114
115 /* The ARB_compute_variable_group_size spec says:
116 *
117 * "An INVALID_OPERATION error is generated by
118 * DispatchComputeGroupSizeARB if the active program for the compute
119 * shader stage has a fixed work group size."
120 */
121 struct gl_program *prog = ctx->_Shader->CurrentProgram[MESA_SHADER_COMPUTE];
122 if (!prog->info.workgroup_size_variable) {
123 _mesa_error(ctx, GL_INVALID_OPERATION,
124 "glDispatchComputeGroupSizeARB(fixed work group size "
125 "forbidden)");
126 return GL_FALSE;
127 }
128
129 for (int i = 0; i < 3; i++) {
130 /* The ARB_compute_variable_group_size spec says:
131 *
132 * "An INVALID_VALUE error is generated if any of num_groups_x,
133 * num_groups_y and num_groups_z are greater than or equal to the
134 * maximum work group count for the corresponding dimension."
135 */
136 if (info->grid[i] > ctx->Const.MaxComputeWorkGroupCount[i]) {
137 _mesa_error(ctx, GL_INVALID_VALUE,
138 "glDispatchComputeGroupSizeARB(num_groups_%c)", 'x' + i);
139 return GL_FALSE;
140 }
141
142 /* The ARB_compute_variable_group_size spec says:
143 *
144 * "An INVALID_VALUE error is generated by DispatchComputeGroupSizeARB if
145 * any of <group_size_x>, <group_size_y>, or <group_size_z> is less than
146 * or equal to zero or greater than the maximum local work group size
147 * for compute shaders with variable group size
148 * (MAX_COMPUTE_VARIABLE_GROUP_SIZE_ARB) in the corresponding
149 * dimension."
150 *
151 * However, the "less than" is a spec bug because they are declared as
152 * unsigned integers.
153 */
154 if (info->block[i] == 0 ||
155 info->block[i] > ctx->Const.MaxComputeVariableGroupSize[i]) {
156 _mesa_error(ctx, GL_INVALID_VALUE,
157 "glDispatchComputeGroupSizeARB(group_size_%c)", 'x' + i);
158 return GL_FALSE;
159 }
160 }
161
162 /* The ARB_compute_variable_group_size spec says:
163 *
164 * "An INVALID_VALUE error is generated by DispatchComputeGroupSizeARB if
165 * the product of <group_size_x>, <group_size_y>, and <group_size_z> exceeds
166 * the implementation-dependent maximum local work group invocation count
167 * for compute shaders with variable group size
168 * (MAX_COMPUTE_VARIABLE_GROUP_INVOCATIONS_ARB)."
169 */
170 uint64_t total_invocations = info->block[0] * info->block[1];
171 if (total_invocations <= UINT32_MAX) {
172 /* Only bother multiplying the third value if total still fits in
173 * 32-bit, since MaxComputeVariableGroupInvocations is also 32-bit.
174 */
175 total_invocations *= info->block[2];
176 }
177 if (total_invocations > ctx->Const.MaxComputeVariableGroupInvocations) {
178 _mesa_error(ctx, GL_INVALID_VALUE,
179 "glDispatchComputeGroupSizeARB(product of local_sizes "
180 "exceeds MAX_COMPUTE_VARIABLE_GROUP_INVOCATIONS_ARB "
181 "(%u * %u * %u > %u))",
182 info->block[0], info->block[1], info->block[2],
183 ctx->Const.MaxComputeVariableGroupInvocations);
184 return GL_FALSE;
185 }
186
187 /* The NV_compute_shader_derivatives spec says:
188 *
189 * "An INVALID_VALUE error is generated by DispatchComputeGroupSizeARB if
190 * the active program for the compute shader stage has a compute shader
191 * using the "derivative_group_quadsNV" layout qualifier and
192 * <group_size_x> or <group_size_y> is not a multiple of two.
193 *
194 * An INVALID_VALUE error is generated by DispatchComputeGroupSizeARB if
195 * the active program for the compute shader stage has a compute shader
196 * using the "derivative_group_linearNV" layout qualifier and the product
197 * of <group_size_x>, <group_size_y>, and <group_size_z> is not a multiple
198 * of four."
199 */
200 if (prog->info.derivative_group == DERIVATIVE_GROUP_QUADS &&
201 ((info->block[0] & 1) || (info->block[1] & 1))) {
202 _mesa_error(ctx, GL_INVALID_VALUE,
203 "glDispatchComputeGroupSizeARB(derivative_group_quadsNV "
204 "requires group_size_x (%d) and group_size_y (%d) to be "
205 "divisble by 2)", info->block[0], info->block[1]);
206 return GL_FALSE;
207 }
208
209 if (prog->info.derivative_group == DERIVATIVE_GROUP_LINEAR &&
210 total_invocations & 3) {
211 _mesa_error(ctx, GL_INVALID_VALUE,
212 "glDispatchComputeGroupSizeARB(derivative_group_linearNV "
213 "requires product of group sizes (%"PRIu64") to be divisible "
214 "by 4)", total_invocations);
215 return GL_FALSE;
216 }
217
218 return GL_TRUE;
219 }
220
221 static bool
valid_dispatch_indirect(struct gl_context * ctx,GLintptr indirect)222 valid_dispatch_indirect(struct gl_context *ctx, GLintptr indirect)
223 {
224 GLsizei size = 3 * sizeof(GLuint);
225 const uint64_t end = (uint64_t) indirect + size;
226 const char *name = "glDispatchComputeIndirect";
227
228 if (!check_valid_to_compute(ctx, name))
229 return GL_FALSE;
230
231 /* From the OpenGL 4.3 Core Specification, Chapter 19, Compute Shaders:
232 *
233 * "An INVALID_VALUE error is generated if indirect is negative or is not a
234 * multiple of four."
235 */
236 if (indirect & (sizeof(GLuint) - 1)) {
237 _mesa_error(ctx, GL_INVALID_VALUE,
238 "%s(indirect is not aligned)", name);
239 return GL_FALSE;
240 }
241
242 if (indirect < 0) {
243 _mesa_error(ctx, GL_INVALID_VALUE,
244 "%s(indirect is less than zero)", name);
245 return GL_FALSE;
246 }
247
248 /* From the OpenGL 4.3 Core Specification, Chapter 19, Compute Shaders:
249 *
250 * "An INVALID_OPERATION error is generated if no buffer is bound to the
251 * DRAW_INDIRECT_BUFFER binding, or if the command would source data
252 * beyond the end of the buffer object."
253 */
254 if (!ctx->DispatchIndirectBuffer) {
255 _mesa_error(ctx, GL_INVALID_OPERATION,
256 "%s: no buffer bound to DISPATCH_INDIRECT_BUFFER", name);
257 return GL_FALSE;
258 }
259
260 if (_mesa_check_disallowed_mapping(ctx->DispatchIndirectBuffer)) {
261 _mesa_error(ctx, GL_INVALID_OPERATION,
262 "%s(DISPATCH_INDIRECT_BUFFER is mapped)", name);
263 return GL_FALSE;
264 }
265
266 if (ctx->DispatchIndirectBuffer->Size < end) {
267 _mesa_error(ctx, GL_INVALID_OPERATION,
268 "%s(DISPATCH_INDIRECT_BUFFER too small)", name);
269 return GL_FALSE;
270 }
271
272 /* The ARB_compute_variable_group_size spec says:
273 *
274 * "An INVALID_OPERATION error is generated if the active program for the
275 * compute shader stage has a variable work group size."
276 */
277 struct gl_program *prog = ctx->_Shader->CurrentProgram[MESA_SHADER_COMPUTE];
278 if (prog->info.workgroup_size_variable) {
279 _mesa_error(ctx, GL_INVALID_OPERATION,
280 "%s(variable work group size forbidden)", name);
281 return GL_FALSE;
282 }
283
284 return GL_TRUE;
285 }
286
287 static void
prepare_compute(struct gl_context * ctx)288 prepare_compute(struct gl_context *ctx)
289 {
290 struct st_context *st = st_context(ctx);
291
292 st_flush_bitmap_cache(st);
293 st_invalidate_readpix_cache(st);
294
295 if (ctx->NewState)
296 _mesa_update_state(ctx);
297
298 st_validate_state(st, ST_PIPELINE_COMPUTE_STATE_MASK);
299 }
300
301 static ALWAYS_INLINE void
dispatch_compute(GLuint num_groups_x,GLuint num_groups_y,GLuint num_groups_z,bool no_error)302 dispatch_compute(GLuint num_groups_x, GLuint num_groups_y,
303 GLuint num_groups_z, bool no_error)
304 {
305 GET_CURRENT_CONTEXT(ctx);
306 struct pipe_grid_info info = { 0 };
307
308 FLUSH_VERTICES(ctx, 0, 0);
309
310 if (MESA_VERBOSE & VERBOSE_API)
311 _mesa_debug(ctx, "glDispatchCompute(%d, %d, %d)\n",
312 num_groups_x, num_groups_y, num_groups_z);
313
314 info.grid[0] = num_groups_x;
315 info.grid[1] = num_groups_y;
316 info.grid[2] = num_groups_z;
317
318 if (!no_error && !validate_DispatchCompute(ctx, &info))
319 return;
320
321 if (num_groups_x == 0u || num_groups_y == 0u || num_groups_z == 0u)
322 return;
323
324 struct gl_program *prog =
325 ctx->_Shader->CurrentProgram[MESA_SHADER_COMPUTE];
326 info.block[0] = prog->info.workgroup_size[0];
327 info.block[1] = prog->info.workgroup_size[1];
328 info.block[2] = prog->info.workgroup_size[2];
329
330 prepare_compute(ctx);
331 ctx->pipe->launch_grid(ctx->pipe, &info);
332
333 if (MESA_DEBUG_FLAGS & DEBUG_ALWAYS_FLUSH)
334 _mesa_flush(ctx);
335 }
336
337 void GLAPIENTRY
_mesa_DispatchCompute_no_error(GLuint num_groups_x,GLuint num_groups_y,GLuint num_groups_z)338 _mesa_DispatchCompute_no_error(GLuint num_groups_x, GLuint num_groups_y,
339 GLuint num_groups_z)
340 {
341 dispatch_compute(num_groups_x, num_groups_y, num_groups_z, true);
342 }
343
344 void GLAPIENTRY
_mesa_DispatchCompute(GLuint num_groups_x,GLuint num_groups_y,GLuint num_groups_z)345 _mesa_DispatchCompute(GLuint num_groups_x,
346 GLuint num_groups_y,
347 GLuint num_groups_z)
348 {
349 dispatch_compute(num_groups_x, num_groups_y, num_groups_z, false);
350 }
351
352 static ALWAYS_INLINE void
dispatch_compute_indirect(GLintptr indirect,bool no_error)353 dispatch_compute_indirect(GLintptr indirect, bool no_error)
354 {
355 GET_CURRENT_CONTEXT(ctx);
356
357 FLUSH_VERTICES(ctx, 0, 0);
358
359 if (MESA_VERBOSE & VERBOSE_API)
360 _mesa_debug(ctx, "glDispatchComputeIndirect(%ld)\n", (long) indirect);
361
362 if (!no_error && !valid_dispatch_indirect(ctx, indirect))
363 return;
364
365 struct pipe_grid_info info = { 0 };
366 info.indirect_offset = indirect;
367 info.indirect = ctx->DispatchIndirectBuffer->buffer;
368
369 struct gl_program *prog =
370 ctx->_Shader->CurrentProgram[MESA_SHADER_COMPUTE];
371 info.block[0] = prog->info.workgroup_size[0];
372 info.block[1] = prog->info.workgroup_size[1];
373 info.block[2] = prog->info.workgroup_size[2];
374
375 prepare_compute(ctx);
376 ctx->pipe->launch_grid(ctx->pipe, &info);
377
378 if (MESA_DEBUG_FLAGS & DEBUG_ALWAYS_FLUSH)
379 _mesa_flush(ctx);
380 }
381
382 extern void GLAPIENTRY
_mesa_DispatchComputeIndirect_no_error(GLintptr indirect)383 _mesa_DispatchComputeIndirect_no_error(GLintptr indirect)
384 {
385 dispatch_compute_indirect(indirect, true);
386 }
387
388 extern void GLAPIENTRY
_mesa_DispatchComputeIndirect(GLintptr indirect)389 _mesa_DispatchComputeIndirect(GLintptr indirect)
390 {
391 dispatch_compute_indirect(indirect, false);
392 }
393
394 static ALWAYS_INLINE void
dispatch_compute_group_size(GLuint num_groups_x,GLuint num_groups_y,GLuint num_groups_z,GLuint group_size_x,GLuint group_size_y,GLuint group_size_z,bool no_error)395 dispatch_compute_group_size(GLuint num_groups_x, GLuint num_groups_y,
396 GLuint num_groups_z, GLuint group_size_x,
397 GLuint group_size_y, GLuint group_size_z,
398 bool no_error)
399 {
400 GET_CURRENT_CONTEXT(ctx);
401 FLUSH_VERTICES(ctx, 0, 0);
402
403 if (MESA_VERBOSE & VERBOSE_API)
404 _mesa_debug(ctx,
405 "glDispatchComputeGroupSizeARB(%d, %d, %d, %d, %d, %d)\n",
406 num_groups_x, num_groups_y, num_groups_z,
407 group_size_x, group_size_y, group_size_z);
408
409 struct pipe_grid_info info = { 0 };
410 info.grid[0] = num_groups_x;
411 info.grid[1] = num_groups_y;
412 info.grid[2] = num_groups_z;
413
414 info.block[0] = group_size_x;
415 info.block[1] = group_size_y;
416 info.block[2] = group_size_z;
417
418 if (!no_error &&
419 !validate_DispatchComputeGroupSizeARB(ctx, &info))
420 return;
421
422 if (num_groups_x == 0u || num_groups_y == 0u || num_groups_z == 0u)
423 return;
424
425 prepare_compute(ctx);
426 ctx->pipe->launch_grid(ctx->pipe, &info);
427
428 if (MESA_DEBUG_FLAGS & DEBUG_ALWAYS_FLUSH)
429 _mesa_flush(ctx);
430 }
431
432 void GLAPIENTRY
_mesa_DispatchComputeGroupSizeARB_no_error(GLuint num_groups_x,GLuint num_groups_y,GLuint num_groups_z,GLuint group_size_x,GLuint group_size_y,GLuint group_size_z)433 _mesa_DispatchComputeGroupSizeARB_no_error(GLuint num_groups_x,
434 GLuint num_groups_y,
435 GLuint num_groups_z,
436 GLuint group_size_x,
437 GLuint group_size_y,
438 GLuint group_size_z)
439 {
440 dispatch_compute_group_size(num_groups_x, num_groups_y, num_groups_z,
441 group_size_x, group_size_y, group_size_z,
442 true);
443 }
444
445 void GLAPIENTRY
_mesa_DispatchComputeGroupSizeARB(GLuint num_groups_x,GLuint num_groups_y,GLuint num_groups_z,GLuint group_size_x,GLuint group_size_y,GLuint group_size_z)446 _mesa_DispatchComputeGroupSizeARB(GLuint num_groups_x, GLuint num_groups_y,
447 GLuint num_groups_z, GLuint group_size_x,
448 GLuint group_size_y, GLuint group_size_z)
449 {
450 dispatch_compute_group_size(num_groups_x, num_groups_y, num_groups_z,
451 group_size_x, group_size_y, group_size_z,
452 false);
453 }
454