1 /**************************************************************************
2 *
3 * Copyright 2009-2010 VMware, Inc.
4 * All Rights Reserved.
5 *
6 * Permission is hereby granted, free of charge, to any person obtaining a
7 * copy of this software and associated documentation files (the
8 * "Software"), to deal in the Software without restriction, including
9 * without limitation the rights to use, copy, modify, merge, publish,
10 * distribute, sub license, and/or sell copies of the Software, and to
11 * permit persons to whom the Software is furnished to do so, subject to
12 * the following conditions:
13 *
14 * The above copyright notice and this permission notice (including the
15 * next paragraph) shall be included in all copies or substantial portions
16 * of the Software.
17 *
18 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS
19 * OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
20 * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NON-INFRINGEMENT.
21 * IN NO EVENT SHALL VMWARE, INC AND/OR ITS SUPPLIERS BE LIABLE FOR
22 * ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
23 * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
24 * SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
25 *
26 **************************************************************************/
27
28
29 #include "pipe/p_screen.h"
30 #include "pipe/p_context.h"
31 #include "pipe/p_state.h"
32 #include "tgsi/tgsi_ureg.h"
33 #include "tgsi/tgsi_build.h"
34 #include "tgsi/tgsi_from_mesa.h"
35 #include "tgsi/tgsi_info.h"
36 #include "tgsi/tgsi_dump.h"
37 #include "tgsi/tgsi_sanity.h"
38 #include "util/glheader.h"
39 #include "util/u_debug.h"
40 #include "util/u_inlines.h"
41 #include "util/u_memory.h"
42 #include "util/u_math.h"
43 #include "util/u_prim.h"
44 #include "util/u_bitmask.h"
45 #include "compiler/shader_info.h"
46
47 union tgsi_any_token {
48 struct tgsi_header header;
49 struct tgsi_processor processor;
50 struct tgsi_token token;
51 struct tgsi_property prop;
52 struct tgsi_property_data prop_data;
53 struct tgsi_declaration decl;
54 struct tgsi_declaration_range decl_range;
55 struct tgsi_declaration_dimension decl_dim;
56 struct tgsi_declaration_interp decl_interp;
57 struct tgsi_declaration_image decl_image;
58 struct tgsi_declaration_semantic decl_semantic;
59 struct tgsi_declaration_sampler_view decl_sampler_view;
60 struct tgsi_declaration_array array;
61 struct tgsi_immediate imm;
62 union tgsi_immediate_data imm_data;
63 struct tgsi_instruction insn;
64 struct tgsi_instruction_label insn_label;
65 struct tgsi_instruction_texture insn_texture;
66 struct tgsi_instruction_memory insn_memory;
67 struct tgsi_texture_offset insn_texture_offset;
68 struct tgsi_src_register src;
69 struct tgsi_ind_register ind;
70 struct tgsi_dimension dim;
71 struct tgsi_dst_register dst;
72 unsigned value;
73 };
74
75
76 struct ureg_tokens {
77 union tgsi_any_token *tokens;
78 unsigned size;
79 unsigned order;
80 unsigned count;
81 };
82
83 #define UREG_MAX_INPUT (4 * PIPE_MAX_SHADER_INPUTS)
84 #define UREG_MAX_SYSTEM_VALUE PIPE_MAX_ATTRIBS
85 #define UREG_MAX_OUTPUT (4 * PIPE_MAX_SHADER_OUTPUTS)
86 #define UREG_MAX_CONSTANT_RANGE 32
87 #define UREG_MAX_HW_ATOMIC_RANGE 32
88 #define UREG_MAX_IMMEDIATE 4096
89 #define UREG_MAX_ADDR 3
90 #define UREG_MAX_ARRAY_TEMPS 256
91
92 struct const_decl {
93 struct {
94 unsigned first;
95 unsigned last;
96 } constant_range[UREG_MAX_CONSTANT_RANGE];
97 unsigned nr_constant_ranges;
98 };
99
100 struct hw_atomic_decl {
101 struct hw_atomic_decl_range {
102 unsigned first;
103 unsigned last;
104 unsigned array_id;
105 } hw_atomic_range[UREG_MAX_HW_ATOMIC_RANGE];
106 unsigned nr_hw_atomic_ranges;
107 };
108
109 #define DOMAIN_DECL 0
110 #define DOMAIN_INSN 1
111
112 struct ureg_program
113 {
114 enum pipe_shader_type processor;
115 bool supports_any_inout_decl_range;
116 int next_shader_processor;
117
118 struct ureg_input_decl {
119 enum tgsi_semantic semantic_name;
120 unsigned semantic_index;
121 enum tgsi_interpolate_mode interp;
122 unsigned char usage_mask;
123 enum tgsi_interpolate_loc interp_location;
124 unsigned first;
125 unsigned last;
126 unsigned array_id;
127 } input[UREG_MAX_INPUT];
128 unsigned nr_inputs, nr_input_regs;
129
130 unsigned vs_inputs[PIPE_MAX_ATTRIBS/32];
131
132 struct {
133 enum tgsi_semantic semantic_name;
134 unsigned semantic_index;
135 } system_value[UREG_MAX_SYSTEM_VALUE];
136 unsigned nr_system_values;
137
138 struct ureg_output_decl {
139 enum tgsi_semantic semantic_name;
140 unsigned semantic_index;
141 unsigned streams;
142 unsigned usage_mask; /* = TGSI_WRITEMASK_* */
143 unsigned first;
144 unsigned last;
145 unsigned array_id;
146 bool invariant;
147 unsigned value_type; /* = TGSI_RETURN_TYPE_* */
148 } output[UREG_MAX_OUTPUT];
149 unsigned nr_outputs, nr_output_regs;
150
151 struct {
152 union {
153 float f[4];
154 unsigned u[4];
155 int i[4];
156 } value;
157 unsigned nr;
158 unsigned type;
159 } immediate[UREG_MAX_IMMEDIATE];
160 unsigned nr_immediates;
161
162 struct ureg_src sampler[PIPE_MAX_SAMPLERS];
163 unsigned nr_samplers;
164
165 struct {
166 unsigned index;
167 enum tgsi_texture_type target;
168 enum tgsi_return_type return_type_x;
169 enum tgsi_return_type return_type_y;
170 enum tgsi_return_type return_type_z;
171 enum tgsi_return_type return_type_w;
172 } sampler_view[PIPE_MAX_SHADER_SAMPLER_VIEWS];
173 unsigned nr_sampler_views;
174
175 struct {
176 unsigned index;
177 enum tgsi_texture_type target;
178 enum pipe_format format;
179 bool wr;
180 bool raw;
181 } image[PIPE_MAX_SHADER_IMAGES];
182 unsigned nr_images;
183
184 struct {
185 unsigned index;
186 bool atomic;
187 } buffer[PIPE_MAX_SHADER_BUFFERS];
188 unsigned nr_buffers;
189
190 struct util_bitmask *free_temps;
191 struct util_bitmask *local_temps;
192 struct util_bitmask *decl_temps;
193 unsigned nr_temps;
194
195 unsigned array_temps[UREG_MAX_ARRAY_TEMPS];
196 unsigned nr_array_temps;
197
198 struct const_decl const_decls[PIPE_MAX_CONSTANT_BUFFERS];
199
200 struct hw_atomic_decl hw_atomic_decls[PIPE_MAX_HW_ATOMIC_BUFFERS];
201
202 unsigned properties[TGSI_PROPERTY_COUNT];
203
204 unsigned nr_addrs;
205 unsigned nr_instructions;
206
207 struct ureg_tokens domain[2];
208
209 bool use_memory[TGSI_MEMORY_TYPE_COUNT];
210
211 bool precise;
212 };
213
214 static union tgsi_any_token error_tokens[32];
215
tokens_error(struct ureg_tokens * tokens)216 static void tokens_error( struct ureg_tokens *tokens )
217 {
218 if (tokens->tokens && tokens->tokens != error_tokens)
219 FREE(tokens->tokens);
220
221 tokens->tokens = error_tokens;
222 tokens->size = ARRAY_SIZE(error_tokens);
223 tokens->count = 0;
224 }
225
226
tokens_expand(struct ureg_tokens * tokens,unsigned count)227 static void tokens_expand( struct ureg_tokens *tokens,
228 unsigned count )
229 {
230 unsigned old_size = tokens->size * sizeof(unsigned);
231
232 if (tokens->tokens == error_tokens) {
233 return;
234 }
235
236 while (tokens->count + count > tokens->size) {
237 tokens->size = (1 << ++tokens->order);
238 }
239
240 tokens->tokens = REALLOC(tokens->tokens,
241 old_size,
242 tokens->size * sizeof(unsigned));
243 if (tokens->tokens == NULL) {
244 tokens_error(tokens);
245 }
246 }
247
set_bad(struct ureg_program * ureg)248 static void set_bad( struct ureg_program *ureg )
249 {
250 tokens_error(&ureg->domain[0]);
251 }
252
253
254
get_tokens(struct ureg_program * ureg,unsigned domain,unsigned count)255 static union tgsi_any_token *get_tokens( struct ureg_program *ureg,
256 unsigned domain,
257 unsigned count )
258 {
259 struct ureg_tokens *tokens = &ureg->domain[domain];
260 union tgsi_any_token *result;
261
262 if (tokens->count + count > tokens->size)
263 tokens_expand(tokens, count);
264
265 result = &tokens->tokens[tokens->count];
266 tokens->count += count;
267 return result;
268 }
269
270
retrieve_token(struct ureg_program * ureg,unsigned domain,unsigned nr)271 static union tgsi_any_token *retrieve_token( struct ureg_program *ureg,
272 unsigned domain,
273 unsigned nr )
274 {
275 if (ureg->domain[domain].tokens == error_tokens)
276 return &error_tokens[0];
277
278 return &ureg->domain[domain].tokens[nr];
279 }
280
281
282 void
ureg_property(struct ureg_program * ureg,unsigned name,unsigned value)283 ureg_property(struct ureg_program *ureg, unsigned name, unsigned value)
284 {
285 assert(name < ARRAY_SIZE(ureg->properties));
286 ureg->properties[name] = value;
287 }
288
289 struct ureg_src
ureg_DECL_fs_input_centroid_layout(struct ureg_program * ureg,enum tgsi_semantic semantic_name,unsigned semantic_index,enum tgsi_interpolate_mode interp_mode,enum tgsi_interpolate_loc interp_location,unsigned index,unsigned usage_mask,unsigned array_id,unsigned array_size)290 ureg_DECL_fs_input_centroid_layout(struct ureg_program *ureg,
291 enum tgsi_semantic semantic_name,
292 unsigned semantic_index,
293 enum tgsi_interpolate_mode interp_mode,
294 enum tgsi_interpolate_loc interp_location,
295 unsigned index,
296 unsigned usage_mask,
297 unsigned array_id,
298 unsigned array_size)
299 {
300 unsigned i;
301
302 assert(usage_mask != 0);
303 assert(usage_mask <= TGSI_WRITEMASK_XYZW);
304
305 for (i = 0; i < ureg->nr_inputs; i++) {
306 if (ureg->input[i].semantic_name == semantic_name &&
307 ureg->input[i].semantic_index == semantic_index) {
308 assert(ureg->input[i].interp == interp_mode);
309 assert(ureg->input[i].interp_location == interp_location);
310 if (ureg->input[i].array_id == array_id) {
311 ureg->input[i].usage_mask |= usage_mask;
312 ureg->input[i].last = MAX2(ureg->input[i].last, ureg->input[i].first + array_size - 1);
313 ureg->nr_input_regs = MAX2(ureg->nr_input_regs, ureg->input[i].last + 1);
314 goto out;
315 }
316 assert((ureg->input[i].usage_mask & usage_mask) == 0);
317 }
318 }
319
320 if (ureg->nr_inputs < UREG_MAX_INPUT) {
321 assert(array_size >= 1);
322 ureg->input[i].semantic_name = semantic_name;
323 ureg->input[i].semantic_index = semantic_index;
324 ureg->input[i].interp = interp_mode;
325 ureg->input[i].interp_location = interp_location;
326 ureg->input[i].first = index;
327 ureg->input[i].last = index + array_size - 1;
328 ureg->input[i].array_id = array_id;
329 ureg->input[i].usage_mask = usage_mask;
330 ureg->nr_input_regs = MAX2(ureg->nr_input_regs, index + array_size);
331 ureg->nr_inputs++;
332 } else {
333 set_bad(ureg);
334 }
335
336 out:
337 return ureg_src_array_register(TGSI_FILE_INPUT, ureg->input[i].first,
338 array_id);
339 }
340
341 struct ureg_src
ureg_DECL_fs_input_centroid(struct ureg_program * ureg,enum tgsi_semantic semantic_name,unsigned semantic_index,enum tgsi_interpolate_mode interp_mode,enum tgsi_interpolate_loc interp_location,unsigned array_id,unsigned array_size)342 ureg_DECL_fs_input_centroid(struct ureg_program *ureg,
343 enum tgsi_semantic semantic_name,
344 unsigned semantic_index,
345 enum tgsi_interpolate_mode interp_mode,
346 enum tgsi_interpolate_loc interp_location,
347 unsigned array_id,
348 unsigned array_size)
349 {
350 return ureg_DECL_fs_input_centroid_layout(ureg,
351 semantic_name, semantic_index, interp_mode,
352 interp_location,
353 ureg->nr_input_regs, TGSI_WRITEMASK_XYZW, array_id, array_size);
354 }
355
356
357 struct ureg_src
ureg_DECL_vs_input(struct ureg_program * ureg,unsigned index)358 ureg_DECL_vs_input( struct ureg_program *ureg,
359 unsigned index )
360 {
361 assert(ureg->processor == PIPE_SHADER_VERTEX);
362 assert(index / 32 < ARRAY_SIZE(ureg->vs_inputs));
363
364 ureg->vs_inputs[index/32] |= 1 << (index % 32);
365 return ureg_src_register( TGSI_FILE_INPUT, index );
366 }
367
368
369 struct ureg_src
ureg_DECL_input_layout(struct ureg_program * ureg,enum tgsi_semantic semantic_name,unsigned semantic_index,unsigned index,unsigned usage_mask,unsigned array_id,unsigned array_size)370 ureg_DECL_input_layout(struct ureg_program *ureg,
371 enum tgsi_semantic semantic_name,
372 unsigned semantic_index,
373 unsigned index,
374 unsigned usage_mask,
375 unsigned array_id,
376 unsigned array_size)
377 {
378 return ureg_DECL_fs_input_centroid_layout(ureg,
379 semantic_name, semantic_index,
380 TGSI_INTERPOLATE_CONSTANT, TGSI_INTERPOLATE_LOC_CENTER,
381 index, usage_mask, array_id, array_size);
382 }
383
384
385 struct ureg_src
ureg_DECL_input(struct ureg_program * ureg,enum tgsi_semantic semantic_name,unsigned semantic_index,unsigned array_id,unsigned array_size)386 ureg_DECL_input(struct ureg_program *ureg,
387 enum tgsi_semantic semantic_name,
388 unsigned semantic_index,
389 unsigned array_id,
390 unsigned array_size)
391 {
392 return ureg_DECL_fs_input_centroid(ureg, semantic_name, semantic_index,
393 TGSI_INTERPOLATE_CONSTANT,
394 TGSI_INTERPOLATE_LOC_CENTER,
395 array_id, array_size);
396 }
397
398
399 struct ureg_src
ureg_DECL_system_value(struct ureg_program * ureg,enum tgsi_semantic semantic_name,unsigned semantic_index)400 ureg_DECL_system_value(struct ureg_program *ureg,
401 enum tgsi_semantic semantic_name,
402 unsigned semantic_index)
403 {
404 unsigned i;
405
406 for (i = 0; i < ureg->nr_system_values; i++) {
407 if (ureg->system_value[i].semantic_name == semantic_name &&
408 ureg->system_value[i].semantic_index == semantic_index) {
409 goto out;
410 }
411 }
412
413 if (ureg->nr_system_values < UREG_MAX_SYSTEM_VALUE) {
414 ureg->system_value[ureg->nr_system_values].semantic_name = semantic_name;
415 ureg->system_value[ureg->nr_system_values].semantic_index = semantic_index;
416 i = ureg->nr_system_values;
417 ureg->nr_system_values++;
418 } else {
419 set_bad(ureg);
420 }
421
422 out:
423 return ureg_src_register(TGSI_FILE_SYSTEM_VALUE, i);
424 }
425
426
427 struct ureg_dst
ureg_DECL_output_layout(struct ureg_program * ureg,enum tgsi_semantic semantic_name,unsigned semantic_index,unsigned streams,unsigned index,unsigned usage_mask,unsigned array_id,unsigned array_size,bool invariant)428 ureg_DECL_output_layout(struct ureg_program *ureg,
429 enum tgsi_semantic semantic_name,
430 unsigned semantic_index,
431 unsigned streams,
432 unsigned index,
433 unsigned usage_mask,
434 unsigned array_id,
435 unsigned array_size,
436 bool invariant)
437 {
438 unsigned i;
439
440 assert(usage_mask != 0);
441 assert(!(streams & 0x03) || (usage_mask & 1));
442 assert(!(streams & 0x0c) || (usage_mask & 2));
443 assert(!(streams & 0x30) || (usage_mask & 4));
444 assert(!(streams & 0xc0) || (usage_mask & 8));
445
446 for (i = 0; i < ureg->nr_outputs; i++) {
447 if (ureg->output[i].semantic_name == semantic_name &&
448 ureg->output[i].semantic_index == semantic_index) {
449 if (ureg->output[i].array_id == array_id) {
450 ureg->output[i].usage_mask |= usage_mask;
451 ureg->output[i].last = MAX2(ureg->output[i].last, ureg->output[i].first + array_size - 1);
452 ureg->nr_output_regs = MAX2(ureg->nr_output_regs, ureg->output[i].last + 1);
453 goto out;
454 }
455 assert((ureg->output[i].usage_mask & usage_mask) == 0);
456 }
457 }
458
459 if (ureg->nr_outputs < UREG_MAX_OUTPUT) {
460 ureg->output[i].semantic_name = semantic_name;
461 ureg->output[i].semantic_index = semantic_index;
462 ureg->output[i].usage_mask = usage_mask;
463 ureg->output[i].first = index;
464 ureg->output[i].last = index + array_size - 1;
465 ureg->output[i].array_id = array_id;
466 ureg->output[i].invariant = invariant;
467 ureg->nr_output_regs = MAX2(ureg->nr_output_regs, index + array_size);
468 ureg->nr_outputs++;
469 }
470 else {
471 set_bad( ureg );
472 i = 0;
473 }
474
475 out:
476 ureg->output[i].streams |= streams;
477
478 return ureg_dst_array_register(TGSI_FILE_OUTPUT, ureg->output[i].first,
479 array_id);
480 }
481
482
483 struct ureg_dst
ureg_DECL_output_masked(struct ureg_program * ureg,enum tgsi_semantic name,unsigned index,unsigned usage_mask,unsigned array_id,unsigned array_size)484 ureg_DECL_output_masked(struct ureg_program *ureg,
485 enum tgsi_semantic name,
486 unsigned index,
487 unsigned usage_mask,
488 unsigned array_id,
489 unsigned array_size)
490 {
491 return ureg_DECL_output_layout(ureg, name, index, 0,
492 ureg->nr_output_regs, usage_mask, array_id,
493 array_size, false);
494 }
495
496
497 struct ureg_dst
ureg_DECL_output(struct ureg_program * ureg,enum tgsi_semantic name,unsigned index)498 ureg_DECL_output(struct ureg_program *ureg,
499 enum tgsi_semantic name,
500 unsigned index)
501 {
502 return ureg_DECL_output_masked(ureg, name, index, TGSI_WRITEMASK_XYZW,
503 0, 1);
504 }
505
506 struct ureg_dst
ureg_DECL_output_typed(struct ureg_program * ureg,enum tgsi_semantic name,unsigned index,enum tgsi_return_type value_type)507 ureg_DECL_output_typed(struct ureg_program *ureg,
508 enum tgsi_semantic name,
509 unsigned index,
510 enum tgsi_return_type value_type)
511 {
512 struct ureg_dst dst = ureg_DECL_output(ureg, name, index);
513 ureg->output[ureg->nr_outputs - 1].value_type = value_type;
514 return dst;
515 }
516
517 struct ureg_dst
ureg_DECL_output_array(struct ureg_program * ureg,enum tgsi_semantic semantic_name,unsigned semantic_index,unsigned array_id,unsigned array_size)518 ureg_DECL_output_array(struct ureg_program *ureg,
519 enum tgsi_semantic semantic_name,
520 unsigned semantic_index,
521 unsigned array_id,
522 unsigned array_size)
523 {
524 return ureg_DECL_output_masked(ureg, semantic_name, semantic_index,
525 TGSI_WRITEMASK_XYZW,
526 array_id, array_size);
527 }
528
529
530 /* Returns a new constant register. Keep track of which have been
531 * referred to so that we can emit decls later.
532 *
533 * Constant operands declared with this function must be addressed
534 * with a two-dimensional index.
535 *
536 * There is nothing in this code to bind this constant to any tracked
537 * value or manage any constant_buffer contents -- that's the
538 * resposibility of the calling code.
539 */
540 void
ureg_DECL_constant2D(struct ureg_program * ureg,unsigned first,unsigned last,unsigned index2D)541 ureg_DECL_constant2D(struct ureg_program *ureg,
542 unsigned first,
543 unsigned last,
544 unsigned index2D)
545 {
546 struct const_decl *decl = &ureg->const_decls[index2D];
547
548 assert(index2D < PIPE_MAX_CONSTANT_BUFFERS);
549
550 if (decl->nr_constant_ranges < UREG_MAX_CONSTANT_RANGE) {
551 unsigned i = decl->nr_constant_ranges++;
552
553 decl->constant_range[i].first = first;
554 decl->constant_range[i].last = last;
555 }
556 }
557
558
559 /* A one-dimensional, deprecated version of ureg_DECL_constant2D().
560 *
561 * Constant operands declared with this function must be addressed
562 * with a one-dimensional index.
563 */
564 struct ureg_src
ureg_DECL_constant(struct ureg_program * ureg,unsigned index)565 ureg_DECL_constant(struct ureg_program *ureg,
566 unsigned index)
567 {
568 struct const_decl *decl = &ureg->const_decls[0];
569 unsigned minconst = index, maxconst = index;
570 unsigned i;
571
572 /* Inside existing range?
573 */
574 for (i = 0; i < decl->nr_constant_ranges; i++) {
575 if (decl->constant_range[i].first <= index &&
576 decl->constant_range[i].last >= index) {
577 goto out;
578 }
579 }
580
581 /* Extend existing range?
582 */
583 for (i = 0; i < decl->nr_constant_ranges; i++) {
584 if (decl->constant_range[i].last == index - 1) {
585 decl->constant_range[i].last = index;
586 goto out;
587 }
588
589 if (decl->constant_range[i].first == index + 1) {
590 decl->constant_range[i].first = index;
591 goto out;
592 }
593
594 minconst = MIN2(minconst, decl->constant_range[i].first);
595 maxconst = MAX2(maxconst, decl->constant_range[i].last);
596 }
597
598 /* Create new range?
599 */
600 if (decl->nr_constant_ranges < UREG_MAX_CONSTANT_RANGE) {
601 i = decl->nr_constant_ranges++;
602 decl->constant_range[i].first = index;
603 decl->constant_range[i].last = index;
604 goto out;
605 }
606
607 /* Collapse all ranges down to one:
608 */
609 i = 0;
610 decl->constant_range[0].first = minconst;
611 decl->constant_range[0].last = maxconst;
612 decl->nr_constant_ranges = 1;
613
614 out:
615 assert(i < decl->nr_constant_ranges);
616 assert(decl->constant_range[i].first <= index);
617 assert(decl->constant_range[i].last >= index);
618
619 struct ureg_src src = ureg_src_register(TGSI_FILE_CONSTANT, index);
620 return ureg_src_dimension(src, 0);
621 }
622
623
624 /* Returns a new hw atomic register. Keep track of which have been
625 * referred to so that we can emit decls later.
626 */
627 void
ureg_DECL_hw_atomic(struct ureg_program * ureg,unsigned first,unsigned last,unsigned buffer_id,unsigned array_id)628 ureg_DECL_hw_atomic(struct ureg_program *ureg,
629 unsigned first,
630 unsigned last,
631 unsigned buffer_id,
632 unsigned array_id)
633 {
634 struct hw_atomic_decl *decl = &ureg->hw_atomic_decls[buffer_id];
635
636 if (decl->nr_hw_atomic_ranges < UREG_MAX_HW_ATOMIC_RANGE) {
637 unsigned i = decl->nr_hw_atomic_ranges++;
638
639 decl->hw_atomic_range[i].first = first;
640 decl->hw_atomic_range[i].last = last;
641 decl->hw_atomic_range[i].array_id = array_id;
642 } else {
643 set_bad(ureg);
644 }
645 }
646
alloc_temporary(struct ureg_program * ureg,bool local)647 static struct ureg_dst alloc_temporary( struct ureg_program *ureg,
648 bool local )
649 {
650 unsigned i;
651
652 /* Look for a released temporary.
653 */
654 for (i = util_bitmask_get_first_index(ureg->free_temps);
655 i != UTIL_BITMASK_INVALID_INDEX;
656 i = util_bitmask_get_next_index(ureg->free_temps, i + 1)) {
657 if (util_bitmask_get(ureg->local_temps, i) == local)
658 break;
659 }
660
661 /* Or allocate a new one.
662 */
663 if (i == UTIL_BITMASK_INVALID_INDEX) {
664 i = ureg->nr_temps++;
665
666 if (local)
667 util_bitmask_set(ureg->local_temps, i);
668
669 /* Start a new declaration when the local flag changes */
670 if (!i || util_bitmask_get(ureg->local_temps, i - 1) != local)
671 util_bitmask_set(ureg->decl_temps, i);
672 }
673
674 util_bitmask_clear(ureg->free_temps, i);
675
676 return ureg_dst_register( TGSI_FILE_TEMPORARY, i );
677 }
678
ureg_DECL_temporary(struct ureg_program * ureg)679 struct ureg_dst ureg_DECL_temporary( struct ureg_program *ureg )
680 {
681 return alloc_temporary(ureg, false);
682 }
683
ureg_DECL_local_temporary(struct ureg_program * ureg)684 struct ureg_dst ureg_DECL_local_temporary( struct ureg_program *ureg )
685 {
686 return alloc_temporary(ureg, true);
687 }
688
ureg_DECL_array_temporary(struct ureg_program * ureg,unsigned size,bool local)689 struct ureg_dst ureg_DECL_array_temporary( struct ureg_program *ureg,
690 unsigned size,
691 bool local )
692 {
693 unsigned i = ureg->nr_temps;
694 struct ureg_dst dst = ureg_dst_register( TGSI_FILE_TEMPORARY, i );
695
696 if (local)
697 util_bitmask_set(ureg->local_temps, i);
698
699 /* Always start a new declaration at the start */
700 util_bitmask_set(ureg->decl_temps, i);
701
702 ureg->nr_temps += size;
703
704 /* and also at the end of the array */
705 util_bitmask_set(ureg->decl_temps, ureg->nr_temps);
706
707 if (ureg->nr_array_temps < UREG_MAX_ARRAY_TEMPS) {
708 ureg->array_temps[ureg->nr_array_temps++] = i;
709 dst.ArrayID = ureg->nr_array_temps;
710 }
711
712 return dst;
713 }
714
ureg_release_temporary(struct ureg_program * ureg,struct ureg_dst tmp)715 void ureg_release_temporary( struct ureg_program *ureg,
716 struct ureg_dst tmp )
717 {
718 if(tmp.File == TGSI_FILE_TEMPORARY)
719 util_bitmask_set(ureg->free_temps, tmp.Index);
720 }
721
722
723 /* Allocate a new address register.
724 */
ureg_DECL_address(struct ureg_program * ureg)725 struct ureg_dst ureg_DECL_address( struct ureg_program *ureg )
726 {
727 if (ureg->nr_addrs < UREG_MAX_ADDR)
728 return ureg_dst_register( TGSI_FILE_ADDRESS, ureg->nr_addrs++ );
729
730 assert( 0 );
731 return ureg_dst_register( TGSI_FILE_ADDRESS, 0 );
732 }
733
734 /* Allocate a new sampler.
735 */
ureg_DECL_sampler(struct ureg_program * ureg,unsigned nr)736 struct ureg_src ureg_DECL_sampler( struct ureg_program *ureg,
737 unsigned nr )
738 {
739 unsigned i;
740
741 for (i = 0; i < ureg->nr_samplers; i++)
742 if (ureg->sampler[i].Index == (int)nr)
743 return ureg->sampler[i];
744
745 if (i < PIPE_MAX_SAMPLERS) {
746 ureg->sampler[i] = ureg_src_register( TGSI_FILE_SAMPLER, nr );
747 ureg->nr_samplers++;
748 return ureg->sampler[i];
749 }
750
751 assert( 0 );
752 return ureg->sampler[0];
753 }
754
755 /*
756 * Allocate a new shader sampler view.
757 */
758 struct ureg_src
ureg_DECL_sampler_view(struct ureg_program * ureg,unsigned index,enum tgsi_texture_type target,enum tgsi_return_type return_type_x,enum tgsi_return_type return_type_y,enum tgsi_return_type return_type_z,enum tgsi_return_type return_type_w)759 ureg_DECL_sampler_view(struct ureg_program *ureg,
760 unsigned index,
761 enum tgsi_texture_type target,
762 enum tgsi_return_type return_type_x,
763 enum tgsi_return_type return_type_y,
764 enum tgsi_return_type return_type_z,
765 enum tgsi_return_type return_type_w)
766 {
767 struct ureg_src reg = ureg_src_register(TGSI_FILE_SAMPLER_VIEW, index);
768 unsigned i;
769
770 for (i = 0; i < ureg->nr_sampler_views; i++) {
771 if (ureg->sampler_view[i].index == index) {
772 return reg;
773 }
774 }
775
776 if (i < PIPE_MAX_SHADER_SAMPLER_VIEWS) {
777 ureg->sampler_view[i].index = index;
778 ureg->sampler_view[i].target = target;
779 ureg->sampler_view[i].return_type_x = return_type_x;
780 ureg->sampler_view[i].return_type_y = return_type_y;
781 ureg->sampler_view[i].return_type_z = return_type_z;
782 ureg->sampler_view[i].return_type_w = return_type_w;
783 ureg->nr_sampler_views++;
784 return reg;
785 }
786
787 assert(0);
788 return reg;
789 }
790
791 /* Allocate a new image.
792 */
793 struct ureg_src
ureg_DECL_image(struct ureg_program * ureg,unsigned index,enum tgsi_texture_type target,enum pipe_format format,bool wr,bool raw)794 ureg_DECL_image(struct ureg_program *ureg,
795 unsigned index,
796 enum tgsi_texture_type target,
797 enum pipe_format format,
798 bool wr,
799 bool raw)
800 {
801 struct ureg_src reg = ureg_src_register(TGSI_FILE_IMAGE, index);
802 unsigned i;
803
804 for (i = 0; i < ureg->nr_images; i++)
805 if (ureg->image[i].index == index)
806 return reg;
807
808 if (i < PIPE_MAX_SHADER_IMAGES) {
809 ureg->image[i].index = index;
810 ureg->image[i].target = target;
811 ureg->image[i].wr = wr;
812 ureg->image[i].raw = raw;
813 ureg->image[i].format = format;
814 ureg->nr_images++;
815 return reg;
816 }
817
818 assert(0);
819 return reg;
820 }
821
822 /* Allocate a new buffer.
823 */
ureg_DECL_buffer(struct ureg_program * ureg,unsigned nr,bool atomic)824 struct ureg_src ureg_DECL_buffer(struct ureg_program *ureg, unsigned nr,
825 bool atomic)
826 {
827 struct ureg_src reg = ureg_src_register(TGSI_FILE_BUFFER, nr);
828 unsigned i;
829
830 for (i = 0; i < ureg->nr_buffers; i++)
831 if (ureg->buffer[i].index == nr)
832 return reg;
833
834 if (i < PIPE_MAX_SHADER_BUFFERS) {
835 ureg->buffer[i].index = nr;
836 ureg->buffer[i].atomic = atomic;
837 ureg->nr_buffers++;
838 return reg;
839 }
840
841 assert(0);
842 return reg;
843 }
844
845 /* Allocate a memory area.
846 */
ureg_DECL_memory(struct ureg_program * ureg,unsigned memory_type)847 struct ureg_src ureg_DECL_memory(struct ureg_program *ureg,
848 unsigned memory_type)
849 {
850 struct ureg_src reg = ureg_src_register(TGSI_FILE_MEMORY, memory_type);
851
852 ureg->use_memory[memory_type] = true;
853 return reg;
854 }
855
856 static int
match_or_expand_immediate64(const unsigned * v,unsigned nr,unsigned * v2,unsigned * pnr2,unsigned * swizzle)857 match_or_expand_immediate64( const unsigned *v,
858 unsigned nr,
859 unsigned *v2,
860 unsigned *pnr2,
861 unsigned *swizzle )
862 {
863 unsigned nr2 = *pnr2;
864 unsigned i, j;
865 *swizzle = 0;
866
867 for (i = 0; i < nr; i += 2) {
868 bool found = false;
869
870 for (j = 0; j < nr2 && !found; j += 2) {
871 if (v[i] == v2[j] && v[i + 1] == v2[j + 1]) {
872 *swizzle |= (j << (i * 2)) | ((j + 1) << ((i + 1) * 2));
873 found = true;
874 }
875 }
876 if (!found) {
877 if ((nr2) >= 4) {
878 return false;
879 }
880
881 v2[nr2] = v[i];
882 v2[nr2 + 1] = v[i + 1];
883
884 *swizzle |= (nr2 << (i * 2)) | ((nr2 + 1) << ((i + 1) * 2));
885 nr2 += 2;
886 }
887 }
888
889 /* Actually expand immediate only when fully succeeded.
890 */
891 *pnr2 = nr2;
892 return true;
893 }
894
895 static int
match_or_expand_immediate(const unsigned * v,int type,unsigned nr,unsigned * v2,unsigned * pnr2,unsigned * swizzle)896 match_or_expand_immediate( const unsigned *v,
897 int type,
898 unsigned nr,
899 unsigned *v2,
900 unsigned *pnr2,
901 unsigned *swizzle )
902 {
903 unsigned nr2 = *pnr2;
904 unsigned i, j;
905
906 if (type == TGSI_IMM_FLOAT64 ||
907 type == TGSI_IMM_UINT64 ||
908 type == TGSI_IMM_INT64)
909 return match_or_expand_immediate64(v, nr, v2, pnr2, swizzle);
910
911 *swizzle = 0;
912
913 for (i = 0; i < nr; i++) {
914 bool found = false;
915
916 for (j = 0; j < nr2 && !found; j++) {
917 if (v[i] == v2[j]) {
918 *swizzle |= j << (i * 2);
919 found = true;
920 }
921 }
922
923 if (!found) {
924 if (nr2 >= 4) {
925 return false;
926 }
927
928 v2[nr2] = v[i];
929 *swizzle |= nr2 << (i * 2);
930 nr2++;
931 }
932 }
933
934 /* Actually expand immediate only when fully succeeded.
935 */
936 *pnr2 = nr2;
937 return true;
938 }
939
940
941 static struct ureg_src
decl_immediate(struct ureg_program * ureg,const unsigned * v,unsigned nr,unsigned type)942 decl_immediate( struct ureg_program *ureg,
943 const unsigned *v,
944 unsigned nr,
945 unsigned type )
946 {
947 unsigned i, j;
948 unsigned swizzle = 0;
949
950 /* Could do a first pass where we examine all existing immediates
951 * without expanding.
952 */
953
954 for (i = 0; i < ureg->nr_immediates; i++) {
955 if (ureg->immediate[i].type != type) {
956 continue;
957 }
958 if (match_or_expand_immediate(v,
959 type,
960 nr,
961 ureg->immediate[i].value.u,
962 &ureg->immediate[i].nr,
963 &swizzle)) {
964 goto out;
965 }
966 }
967
968 if (ureg->nr_immediates < UREG_MAX_IMMEDIATE) {
969 i = ureg->nr_immediates++;
970 ureg->immediate[i].type = type;
971 if (match_or_expand_immediate(v,
972 type,
973 nr,
974 ureg->immediate[i].value.u,
975 &ureg->immediate[i].nr,
976 &swizzle)) {
977 goto out;
978 }
979 }
980
981 set_bad(ureg);
982
983 out:
984 /* Make sure that all referenced elements are from this immediate.
985 * Has the effect of making size-one immediates into scalars.
986 */
987 if (type == TGSI_IMM_FLOAT64 ||
988 type == TGSI_IMM_UINT64 ||
989 type == TGSI_IMM_INT64) {
990 for (j = nr; j < 4; j+=2) {
991 swizzle |= (swizzle & 0xf) << (j * 2);
992 }
993 } else {
994 for (j = nr; j < 4; j++) {
995 swizzle |= (swizzle & 0x3) << (j * 2);
996 }
997 }
998 return ureg_swizzle(ureg_src_register(TGSI_FILE_IMMEDIATE, i),
999 (swizzle >> 0) & 0x3,
1000 (swizzle >> 2) & 0x3,
1001 (swizzle >> 4) & 0x3,
1002 (swizzle >> 6) & 0x3);
1003 }
1004
1005
1006 struct ureg_src
ureg_DECL_immediate(struct ureg_program * ureg,const float * v,unsigned nr)1007 ureg_DECL_immediate( struct ureg_program *ureg,
1008 const float *v,
1009 unsigned nr )
1010 {
1011 union {
1012 float f[4];
1013 unsigned u[4];
1014 } fu;
1015 unsigned int i;
1016
1017 for (i = 0; i < nr; i++) {
1018 fu.f[i] = v[i];
1019 }
1020
1021 return decl_immediate(ureg, fu.u, nr, TGSI_IMM_FLOAT32);
1022 }
1023
1024 struct ureg_src
ureg_DECL_immediate_f64(struct ureg_program * ureg,const double * v,unsigned nr)1025 ureg_DECL_immediate_f64( struct ureg_program *ureg,
1026 const double *v,
1027 unsigned nr )
1028 {
1029 union {
1030 unsigned u[4];
1031 double d[2];
1032 } fu;
1033 unsigned int i;
1034
1035 assert((nr / 2) < 3);
1036 for (i = 0; i < nr / 2; i++) {
1037 fu.d[i] = v[i];
1038 }
1039
1040 return decl_immediate(ureg, fu.u, nr, TGSI_IMM_FLOAT64);
1041 }
1042
1043 struct ureg_src
ureg_DECL_immediate_uint(struct ureg_program * ureg,const unsigned * v,unsigned nr)1044 ureg_DECL_immediate_uint( struct ureg_program *ureg,
1045 const unsigned *v,
1046 unsigned nr )
1047 {
1048 return decl_immediate(ureg, v, nr, TGSI_IMM_UINT32);
1049 }
1050
1051
1052 struct ureg_src
ureg_DECL_immediate_block_uint(struct ureg_program * ureg,const unsigned * v,unsigned nr)1053 ureg_DECL_immediate_block_uint( struct ureg_program *ureg,
1054 const unsigned *v,
1055 unsigned nr )
1056 {
1057 unsigned index;
1058 unsigned i;
1059
1060 if (ureg->nr_immediates + (nr + 3) / 4 > UREG_MAX_IMMEDIATE) {
1061 set_bad(ureg);
1062 return ureg_src_register(TGSI_FILE_IMMEDIATE, 0);
1063 }
1064
1065 index = ureg->nr_immediates;
1066 ureg->nr_immediates += (nr + 3) / 4;
1067
1068 for (i = index; i < ureg->nr_immediates; i++) {
1069 ureg->immediate[i].type = TGSI_IMM_UINT32;
1070 ureg->immediate[i].nr = nr > 4 ? 4 : nr;
1071 memcpy(ureg->immediate[i].value.u,
1072 &v[(i - index) * 4],
1073 ureg->immediate[i].nr * sizeof(unsigned));
1074 nr -= 4;
1075 }
1076
1077 return ureg_src_register(TGSI_FILE_IMMEDIATE, index);
1078 }
1079
1080
1081 struct ureg_src
ureg_DECL_immediate_int(struct ureg_program * ureg,const int * v,unsigned nr)1082 ureg_DECL_immediate_int( struct ureg_program *ureg,
1083 const int *v,
1084 unsigned nr )
1085 {
1086 return decl_immediate(ureg, (const unsigned *)v, nr, TGSI_IMM_INT32);
1087 }
1088
1089 struct ureg_src
ureg_DECL_immediate_uint64(struct ureg_program * ureg,const uint64_t * v,unsigned nr)1090 ureg_DECL_immediate_uint64( struct ureg_program *ureg,
1091 const uint64_t *v,
1092 unsigned nr )
1093 {
1094 union {
1095 unsigned u[4];
1096 uint64_t u64[2];
1097 } fu;
1098 unsigned int i;
1099
1100 assert((nr / 2) < 3);
1101 for (i = 0; i < nr / 2; i++) {
1102 fu.u64[i] = v[i];
1103 }
1104
1105 return decl_immediate(ureg, fu.u, nr, TGSI_IMM_UINT64);
1106 }
1107
1108 struct ureg_src
ureg_DECL_immediate_int64(struct ureg_program * ureg,const int64_t * v,unsigned nr)1109 ureg_DECL_immediate_int64( struct ureg_program *ureg,
1110 const int64_t *v,
1111 unsigned nr )
1112 {
1113 union {
1114 unsigned u[4];
1115 int64_t i64[2];
1116 } fu;
1117 unsigned int i;
1118
1119 assert((nr / 2) < 3);
1120 for (i = 0; i < nr / 2; i++) {
1121 fu.i64[i] = v[i];
1122 }
1123
1124 return decl_immediate(ureg, fu.u, nr, TGSI_IMM_INT64);
1125 }
1126
1127 void
ureg_emit_src(struct ureg_program * ureg,struct ureg_src src)1128 ureg_emit_src( struct ureg_program *ureg,
1129 struct ureg_src src )
1130 {
1131 unsigned size = 1 + (src.Indirect ? 1 : 0) +
1132 (src.Dimension ? (src.DimIndirect ? 2 : 1) : 0);
1133
1134 union tgsi_any_token *out = get_tokens( ureg, DOMAIN_INSN, size );
1135 unsigned n = 0;
1136
1137 assert(src.File != TGSI_FILE_NULL);
1138 assert(src.File < TGSI_FILE_COUNT);
1139
1140 out[n].value = 0;
1141 out[n].src.File = src.File;
1142 out[n].src.SwizzleX = src.SwizzleX;
1143 out[n].src.SwizzleY = src.SwizzleY;
1144 out[n].src.SwizzleZ = src.SwizzleZ;
1145 out[n].src.SwizzleW = src.SwizzleW;
1146 out[n].src.Index = src.Index;
1147 out[n].src.Negate = src.Negate;
1148 out[0].src.Absolute = src.Absolute;
1149 n++;
1150
1151 if (src.Indirect) {
1152 out[0].src.Indirect = 1;
1153 out[n].value = 0;
1154 out[n].ind.File = src.IndirectFile;
1155 out[n].ind.Swizzle = src.IndirectSwizzle;
1156 out[n].ind.Index = src.IndirectIndex;
1157 if (!ureg->supports_any_inout_decl_range &&
1158 (src.File == TGSI_FILE_INPUT || src.File == TGSI_FILE_OUTPUT))
1159 out[n].ind.ArrayID = 0;
1160 else
1161 out[n].ind.ArrayID = src.ArrayID;
1162 n++;
1163 }
1164
1165 if (src.Dimension) {
1166 out[0].src.Dimension = 1;
1167 out[n].dim.Dimension = 0;
1168 out[n].dim.Padding = 0;
1169 if (src.DimIndirect) {
1170 out[n].dim.Indirect = 1;
1171 out[n].dim.Index = src.DimensionIndex;
1172 n++;
1173 out[n].value = 0;
1174 out[n].ind.File = src.DimIndFile;
1175 out[n].ind.Swizzle = src.DimIndSwizzle;
1176 out[n].ind.Index = src.DimIndIndex;
1177 if (!ureg->supports_any_inout_decl_range &&
1178 (src.File == TGSI_FILE_INPUT || src.File == TGSI_FILE_OUTPUT))
1179 out[n].ind.ArrayID = 0;
1180 else
1181 out[n].ind.ArrayID = src.ArrayID;
1182 } else {
1183 out[n].dim.Indirect = 0;
1184 out[n].dim.Index = src.DimensionIndex;
1185 }
1186 n++;
1187 }
1188
1189 assert(n == size);
1190 }
1191
1192
1193 void
ureg_emit_dst(struct ureg_program * ureg,struct ureg_dst dst)1194 ureg_emit_dst( struct ureg_program *ureg,
1195 struct ureg_dst dst )
1196 {
1197 unsigned size = 1 + (dst.Indirect ? 1 : 0) +
1198 (dst.Dimension ? (dst.DimIndirect ? 2 : 1) : 0);
1199
1200 union tgsi_any_token *out = get_tokens( ureg, DOMAIN_INSN, size );
1201 unsigned n = 0;
1202
1203 assert(dst.File != TGSI_FILE_NULL);
1204 assert(dst.File != TGSI_FILE_SAMPLER);
1205 assert(dst.File != TGSI_FILE_SAMPLER_VIEW);
1206 assert(dst.File != TGSI_FILE_IMMEDIATE);
1207 assert(dst.File < TGSI_FILE_COUNT);
1208
1209 out[n].value = 0;
1210 out[n].dst.File = dst.File;
1211 out[n].dst.WriteMask = dst.WriteMask;
1212 out[n].dst.Indirect = dst.Indirect;
1213 out[n].dst.Index = dst.Index;
1214 n++;
1215
1216 if (dst.Indirect) {
1217 out[n].value = 0;
1218 out[n].ind.File = dst.IndirectFile;
1219 out[n].ind.Swizzle = dst.IndirectSwizzle;
1220 out[n].ind.Index = dst.IndirectIndex;
1221 if (!ureg->supports_any_inout_decl_range &&
1222 (dst.File == TGSI_FILE_INPUT || dst.File == TGSI_FILE_OUTPUT))
1223 out[n].ind.ArrayID = 0;
1224 else
1225 out[n].ind.ArrayID = dst.ArrayID;
1226 n++;
1227 }
1228
1229 if (dst.Dimension) {
1230 out[0].dst.Dimension = 1;
1231 out[n].dim.Dimension = 0;
1232 out[n].dim.Padding = 0;
1233 if (dst.DimIndirect) {
1234 out[n].dim.Indirect = 1;
1235 out[n].dim.Index = dst.DimensionIndex;
1236 n++;
1237 out[n].value = 0;
1238 out[n].ind.File = dst.DimIndFile;
1239 out[n].ind.Swizzle = dst.DimIndSwizzle;
1240 out[n].ind.Index = dst.DimIndIndex;
1241 if (!ureg->supports_any_inout_decl_range &&
1242 (dst.File == TGSI_FILE_INPUT || dst.File == TGSI_FILE_OUTPUT))
1243 out[n].ind.ArrayID = 0;
1244 else
1245 out[n].ind.ArrayID = dst.ArrayID;
1246 } else {
1247 out[n].dim.Indirect = 0;
1248 out[n].dim.Index = dst.DimensionIndex;
1249 }
1250 n++;
1251 }
1252
1253 assert(n == size);
1254 }
1255
1256
validate(enum tgsi_opcode opcode,unsigned nr_dst,unsigned nr_src)1257 static void validate( enum tgsi_opcode opcode,
1258 unsigned nr_dst,
1259 unsigned nr_src )
1260 {
1261 #ifndef NDEBUG
1262 const struct tgsi_opcode_info *info = tgsi_get_opcode_info( opcode );
1263 assert(info);
1264 if (info) {
1265 assert(nr_dst == info->num_dst);
1266 assert(nr_src == info->num_src);
1267 }
1268 #endif
1269 }
1270
1271 struct ureg_emit_insn_result
ureg_emit_insn(struct ureg_program * ureg,enum tgsi_opcode opcode,bool saturate,unsigned precise,unsigned num_dst,unsigned num_src)1272 ureg_emit_insn(struct ureg_program *ureg,
1273 enum tgsi_opcode opcode,
1274 bool saturate,
1275 unsigned precise,
1276 unsigned num_dst,
1277 unsigned num_src)
1278 {
1279 union tgsi_any_token *out;
1280 unsigned count = 1;
1281 struct ureg_emit_insn_result result;
1282
1283 validate( opcode, num_dst, num_src );
1284
1285 out = get_tokens( ureg, DOMAIN_INSN, count );
1286 out[0].insn = tgsi_default_instruction();
1287 out[0].insn.Opcode = opcode;
1288 out[0].insn.Saturate = saturate;
1289 out[0].insn.Precise = precise || ureg->precise;
1290 out[0].insn.NumDstRegs = num_dst;
1291 out[0].insn.NumSrcRegs = num_src;
1292
1293 result.insn_token = ureg->domain[DOMAIN_INSN].count - count;
1294 result.extended_token = result.insn_token;
1295
1296 ureg->nr_instructions++;
1297
1298 return result;
1299 }
1300
1301
1302 /**
1303 * Emit a label token.
1304 * \param label_token returns a token number indicating where the label
1305 * needs to be patched later. Later, this value should be passed to the
1306 * ureg_fixup_label() function.
1307 */
1308 void
ureg_emit_label(struct ureg_program * ureg,unsigned extended_token,unsigned * label_token)1309 ureg_emit_label(struct ureg_program *ureg,
1310 unsigned extended_token,
1311 unsigned *label_token )
1312 {
1313 union tgsi_any_token *out, *insn;
1314
1315 if (!label_token)
1316 return;
1317
1318 out = get_tokens( ureg, DOMAIN_INSN, 1 );
1319 out[0].value = 0;
1320
1321 insn = retrieve_token( ureg, DOMAIN_INSN, extended_token );
1322 insn->insn.Label = 1;
1323
1324 *label_token = ureg->domain[DOMAIN_INSN].count - 1;
1325 }
1326
1327 /* Will return a number which can be used in a label to point to the
1328 * next instruction to be emitted.
1329 */
1330 unsigned
ureg_get_instruction_number(struct ureg_program * ureg)1331 ureg_get_instruction_number( struct ureg_program *ureg )
1332 {
1333 return ureg->nr_instructions;
1334 }
1335
1336 /* Patch a given label (expressed as a token number) to point to a
1337 * given instruction (expressed as an instruction number).
1338 */
1339 void
ureg_fixup_label(struct ureg_program * ureg,unsigned label_token,unsigned instruction_number)1340 ureg_fixup_label(struct ureg_program *ureg,
1341 unsigned label_token,
1342 unsigned instruction_number )
1343 {
1344 union tgsi_any_token *out = retrieve_token( ureg, DOMAIN_INSN, label_token );
1345
1346 out->insn_label.Label = instruction_number;
1347 }
1348
1349
1350 void
ureg_emit_texture(struct ureg_program * ureg,unsigned extended_token,enum tgsi_texture_type target,enum tgsi_return_type return_type,unsigned num_offsets)1351 ureg_emit_texture(struct ureg_program *ureg,
1352 unsigned extended_token,
1353 enum tgsi_texture_type target,
1354 enum tgsi_return_type return_type, unsigned num_offsets)
1355 {
1356 union tgsi_any_token *out, *insn;
1357
1358 out = get_tokens( ureg, DOMAIN_INSN, 1 );
1359 insn = retrieve_token( ureg, DOMAIN_INSN, extended_token );
1360
1361 insn->insn.Texture = 1;
1362
1363 out[0].value = 0;
1364 out[0].insn_texture.Texture = target;
1365 out[0].insn_texture.NumOffsets = num_offsets;
1366 out[0].insn_texture.ReturnType = return_type;
1367 }
1368
1369 void
ureg_emit_texture_offset(struct ureg_program * ureg,const struct tgsi_texture_offset * offset)1370 ureg_emit_texture_offset(struct ureg_program *ureg,
1371 const struct tgsi_texture_offset *offset)
1372 {
1373 union tgsi_any_token *out;
1374
1375 out = get_tokens( ureg, DOMAIN_INSN, 1);
1376
1377 out[0].value = 0;
1378 out[0].insn_texture_offset = *offset;
1379 }
1380
1381 void
ureg_emit_memory(struct ureg_program * ureg,unsigned extended_token,unsigned qualifier,enum tgsi_texture_type texture,enum pipe_format format)1382 ureg_emit_memory(struct ureg_program *ureg,
1383 unsigned extended_token,
1384 unsigned qualifier,
1385 enum tgsi_texture_type texture,
1386 enum pipe_format format)
1387 {
1388 union tgsi_any_token *out, *insn;
1389
1390 out = get_tokens( ureg, DOMAIN_INSN, 1 );
1391 insn = retrieve_token( ureg, DOMAIN_INSN, extended_token );
1392
1393 insn->insn.Memory = 1;
1394
1395 out[0].value = 0;
1396 out[0].insn_memory.Qualifier = qualifier;
1397 out[0].insn_memory.Texture = texture;
1398 out[0].insn_memory.Format = format;
1399 }
1400
1401 void
ureg_fixup_insn_size(struct ureg_program * ureg,unsigned insn)1402 ureg_fixup_insn_size(struct ureg_program *ureg,
1403 unsigned insn )
1404 {
1405 union tgsi_any_token *out = retrieve_token( ureg, DOMAIN_INSN, insn );
1406
1407 assert(out->insn.Type == TGSI_TOKEN_TYPE_INSTRUCTION);
1408 out->insn.NrTokens = ureg->domain[DOMAIN_INSN].count - insn - 1;
1409 }
1410
1411
1412 void
ureg_insn(struct ureg_program * ureg,enum tgsi_opcode opcode,const struct ureg_dst * dst,unsigned nr_dst,const struct ureg_src * src,unsigned nr_src,unsigned precise)1413 ureg_insn(struct ureg_program *ureg,
1414 enum tgsi_opcode opcode,
1415 const struct ureg_dst *dst,
1416 unsigned nr_dst,
1417 const struct ureg_src *src,
1418 unsigned nr_src,
1419 unsigned precise )
1420 {
1421 struct ureg_emit_insn_result insn;
1422 unsigned i;
1423 bool saturate;
1424
1425 if (nr_dst && ureg_dst_is_empty(dst[0])) {
1426 return;
1427 }
1428
1429 saturate = nr_dst ? dst[0].Saturate : false;
1430
1431 insn = ureg_emit_insn(ureg,
1432 opcode,
1433 saturate,
1434 precise,
1435 nr_dst,
1436 nr_src);
1437
1438 for (i = 0; i < nr_dst; i++)
1439 ureg_emit_dst( ureg, dst[i] );
1440
1441 for (i = 0; i < nr_src; i++)
1442 ureg_emit_src( ureg, src[i] );
1443
1444 ureg_fixup_insn_size( ureg, insn.insn_token );
1445 }
1446
1447 void
ureg_tex_insn(struct ureg_program * ureg,enum tgsi_opcode opcode,const struct ureg_dst * dst,unsigned nr_dst,enum tgsi_texture_type target,enum tgsi_return_type return_type,const struct tgsi_texture_offset * texoffsets,unsigned nr_offset,const struct ureg_src * src,unsigned nr_src)1448 ureg_tex_insn(struct ureg_program *ureg,
1449 enum tgsi_opcode opcode,
1450 const struct ureg_dst *dst,
1451 unsigned nr_dst,
1452 enum tgsi_texture_type target,
1453 enum tgsi_return_type return_type,
1454 const struct tgsi_texture_offset *texoffsets,
1455 unsigned nr_offset,
1456 const struct ureg_src *src,
1457 unsigned nr_src )
1458 {
1459 struct ureg_emit_insn_result insn;
1460 unsigned i;
1461 bool saturate;
1462
1463 if (nr_dst && ureg_dst_is_empty(dst[0])) {
1464 return;
1465 }
1466
1467 saturate = nr_dst ? dst[0].Saturate : false;
1468
1469 insn = ureg_emit_insn(ureg,
1470 opcode,
1471 saturate,
1472 0,
1473 nr_dst,
1474 nr_src);
1475
1476 ureg_emit_texture( ureg, insn.extended_token, target, return_type,
1477 nr_offset );
1478
1479 for (i = 0; i < nr_offset; i++)
1480 ureg_emit_texture_offset( ureg, &texoffsets[i]);
1481
1482 for (i = 0; i < nr_dst; i++)
1483 ureg_emit_dst( ureg, dst[i] );
1484
1485 for (i = 0; i < nr_src; i++)
1486 ureg_emit_src( ureg, src[i] );
1487
1488 ureg_fixup_insn_size( ureg, insn.insn_token );
1489 }
1490
1491
1492 void
ureg_memory_insn(struct ureg_program * ureg,enum tgsi_opcode opcode,const struct ureg_dst * dst,unsigned nr_dst,const struct ureg_src * src,unsigned nr_src,unsigned qualifier,enum tgsi_texture_type texture,enum pipe_format format)1493 ureg_memory_insn(struct ureg_program *ureg,
1494 enum tgsi_opcode opcode,
1495 const struct ureg_dst *dst,
1496 unsigned nr_dst,
1497 const struct ureg_src *src,
1498 unsigned nr_src,
1499 unsigned qualifier,
1500 enum tgsi_texture_type texture,
1501 enum pipe_format format)
1502 {
1503 struct ureg_emit_insn_result insn;
1504 unsigned i;
1505
1506 insn = ureg_emit_insn(ureg,
1507 opcode,
1508 false,
1509 0,
1510 nr_dst,
1511 nr_src);
1512
1513 ureg_emit_memory(ureg, insn.extended_token, qualifier, texture, format);
1514
1515 for (i = 0; i < nr_dst; i++)
1516 ureg_emit_dst(ureg, dst[i]);
1517
1518 for (i = 0; i < nr_src; i++)
1519 ureg_emit_src(ureg, src[i]);
1520
1521 ureg_fixup_insn_size(ureg, insn.insn_token);
1522 }
1523
1524
1525 static void
emit_decl_semantic_typed(struct ureg_program * ureg,unsigned file,unsigned first,unsigned last,enum tgsi_semantic semantic_name,unsigned semantic_index,unsigned streams,unsigned usage_mask,unsigned array_id,bool invariant,enum tgsi_return_type value_type)1526 emit_decl_semantic_typed(struct ureg_program *ureg,
1527 unsigned file,
1528 unsigned first,
1529 unsigned last,
1530 enum tgsi_semantic semantic_name,
1531 unsigned semantic_index,
1532 unsigned streams,
1533 unsigned usage_mask,
1534 unsigned array_id,
1535 bool invariant,
1536 enum tgsi_return_type value_type)
1537 {
1538 union tgsi_any_token *out = get_tokens(ureg, DOMAIN_DECL, array_id ? 4 : 3);
1539
1540 out[0].value = 0;
1541 out[0].decl.Type = TGSI_TOKEN_TYPE_DECLARATION;
1542 out[0].decl.NrTokens = 3;
1543 out[0].decl.File = file;
1544 out[0].decl.UsageMask = usage_mask;
1545 out[0].decl.Semantic = 1;
1546 out[0].decl.Array = array_id != 0;
1547 out[0].decl.Invariant = invariant;
1548 out[0].decl.ValueType = value_type;
1549
1550 out[1].value = 0;
1551 out[1].decl_range.First = first;
1552 out[1].decl_range.Last = last;
1553
1554 out[2].value = 0;
1555 out[2].decl_semantic.Name = semantic_name;
1556 out[2].decl_semantic.Index = semantic_index;
1557 out[2].decl_semantic.StreamX = streams & 3;
1558 out[2].decl_semantic.StreamY = (streams >> 2) & 3;
1559 out[2].decl_semantic.StreamZ = (streams >> 4) & 3;
1560 out[2].decl_semantic.StreamW = (streams >> 6) & 3;
1561
1562 if (array_id) {
1563 out[3].value = 0;
1564 out[3].array.ArrayID = array_id;
1565 }
1566 }
1567
1568 static void
emit_decl_semantic(struct ureg_program * ureg,unsigned file,unsigned first,unsigned last,enum tgsi_semantic semantic_name,unsigned semantic_index,unsigned streams,unsigned usage_mask,unsigned array_id,bool invariant)1569 emit_decl_semantic(struct ureg_program *ureg,
1570 unsigned file,
1571 unsigned first,
1572 unsigned last,
1573 enum tgsi_semantic semantic_name,
1574 unsigned semantic_index,
1575 unsigned streams,
1576 unsigned usage_mask,
1577 unsigned array_id,
1578 bool invariant)
1579 {
1580 emit_decl_semantic_typed(ureg, file, first, last,
1581 semantic_name, semantic_index, streams, usage_mask, array_id,
1582 invariant, TGSI_RETURN_TYPE_UNKNOWN);
1583 }
1584
1585 static void
emit_decl_atomic_2d(struct ureg_program * ureg,unsigned first,unsigned last,unsigned index2D,unsigned array_id)1586 emit_decl_atomic_2d(struct ureg_program *ureg,
1587 unsigned first,
1588 unsigned last,
1589 unsigned index2D,
1590 unsigned array_id)
1591 {
1592 union tgsi_any_token *out = get_tokens(ureg, DOMAIN_DECL, array_id ? 4 : 3);
1593
1594 out[0].value = 0;
1595 out[0].decl.Type = TGSI_TOKEN_TYPE_DECLARATION;
1596 out[0].decl.NrTokens = 3;
1597 out[0].decl.File = TGSI_FILE_HW_ATOMIC;
1598 out[0].decl.UsageMask = TGSI_WRITEMASK_XYZW;
1599 out[0].decl.Dimension = 1;
1600 out[0].decl.Array = array_id != 0;
1601
1602 out[1].value = 0;
1603 out[1].decl_range.First = first;
1604 out[1].decl_range.Last = last;
1605
1606 out[2].value = 0;
1607 out[2].decl_dim.Index2D = index2D;
1608
1609 if (array_id) {
1610 out[3].value = 0;
1611 out[3].array.ArrayID = array_id;
1612 }
1613 }
1614
1615 static void
emit_decl_fs(struct ureg_program * ureg,unsigned file,unsigned first,unsigned last,enum tgsi_semantic semantic_name,unsigned semantic_index,enum tgsi_interpolate_mode interpolate,enum tgsi_interpolate_loc interpolate_location,unsigned array_id,unsigned usage_mask)1616 emit_decl_fs(struct ureg_program *ureg,
1617 unsigned file,
1618 unsigned first,
1619 unsigned last,
1620 enum tgsi_semantic semantic_name,
1621 unsigned semantic_index,
1622 enum tgsi_interpolate_mode interpolate,
1623 enum tgsi_interpolate_loc interpolate_location,
1624 unsigned array_id,
1625 unsigned usage_mask)
1626 {
1627 union tgsi_any_token *out = get_tokens(ureg, DOMAIN_DECL,
1628 array_id ? 5 : 4);
1629
1630 out[0].value = 0;
1631 out[0].decl.Type = TGSI_TOKEN_TYPE_DECLARATION;
1632 out[0].decl.NrTokens = 4;
1633 out[0].decl.File = file;
1634 out[0].decl.UsageMask = usage_mask;
1635 out[0].decl.Interpolate = 1;
1636 out[0].decl.Semantic = 1;
1637 out[0].decl.Array = array_id != 0;
1638
1639 out[1].value = 0;
1640 out[1].decl_range.First = first;
1641 out[1].decl_range.Last = last;
1642
1643 out[2].value = 0;
1644 out[2].decl_interp.Interpolate = interpolate;
1645 out[2].decl_interp.Location = interpolate_location;
1646
1647 out[3].value = 0;
1648 out[3].decl_semantic.Name = semantic_name;
1649 out[3].decl_semantic.Index = semantic_index;
1650
1651 if (array_id) {
1652 out[4].value = 0;
1653 out[4].array.ArrayID = array_id;
1654 }
1655 }
1656
1657 static void
emit_decl_temps(struct ureg_program * ureg,unsigned first,unsigned last,bool local,unsigned arrayid)1658 emit_decl_temps( struct ureg_program *ureg,
1659 unsigned first, unsigned last,
1660 bool local,
1661 unsigned arrayid )
1662 {
1663 union tgsi_any_token *out = get_tokens( ureg, DOMAIN_DECL,
1664 arrayid ? 3 : 2 );
1665
1666 out[0].value = 0;
1667 out[0].decl.Type = TGSI_TOKEN_TYPE_DECLARATION;
1668 out[0].decl.NrTokens = 2;
1669 out[0].decl.File = TGSI_FILE_TEMPORARY;
1670 out[0].decl.UsageMask = TGSI_WRITEMASK_XYZW;
1671 out[0].decl.Local = local;
1672
1673 out[1].value = 0;
1674 out[1].decl_range.First = first;
1675 out[1].decl_range.Last = last;
1676
1677 if (arrayid) {
1678 out[0].decl.Array = 1;
1679 out[2].value = 0;
1680 out[2].array.ArrayID = arrayid;
1681 }
1682 }
1683
emit_decl_range(struct ureg_program * ureg,unsigned file,unsigned first,unsigned count)1684 static void emit_decl_range( struct ureg_program *ureg,
1685 unsigned file,
1686 unsigned first,
1687 unsigned count )
1688 {
1689 union tgsi_any_token *out = get_tokens( ureg, DOMAIN_DECL, 2 );
1690
1691 out[0].value = 0;
1692 out[0].decl.Type = TGSI_TOKEN_TYPE_DECLARATION;
1693 out[0].decl.NrTokens = 2;
1694 out[0].decl.File = file;
1695 out[0].decl.UsageMask = TGSI_WRITEMASK_XYZW;
1696 out[0].decl.Semantic = 0;
1697
1698 out[1].value = 0;
1699 out[1].decl_range.First = first;
1700 out[1].decl_range.Last = first + count - 1;
1701 }
1702
1703 static void
emit_decl_range2D(struct ureg_program * ureg,unsigned file,unsigned first,unsigned last,unsigned index2D)1704 emit_decl_range2D(struct ureg_program *ureg,
1705 unsigned file,
1706 unsigned first,
1707 unsigned last,
1708 unsigned index2D)
1709 {
1710 union tgsi_any_token *out = get_tokens(ureg, DOMAIN_DECL, 3);
1711
1712 out[0].value = 0;
1713 out[0].decl.Type = TGSI_TOKEN_TYPE_DECLARATION;
1714 out[0].decl.NrTokens = 3;
1715 out[0].decl.File = file;
1716 out[0].decl.UsageMask = TGSI_WRITEMASK_XYZW;
1717 out[0].decl.Dimension = 1;
1718
1719 out[1].value = 0;
1720 out[1].decl_range.First = first;
1721 out[1].decl_range.Last = last;
1722
1723 out[2].value = 0;
1724 out[2].decl_dim.Index2D = index2D;
1725 }
1726
1727 static void
emit_decl_sampler_view(struct ureg_program * ureg,unsigned index,enum tgsi_texture_type target,enum tgsi_return_type return_type_x,enum tgsi_return_type return_type_y,enum tgsi_return_type return_type_z,enum tgsi_return_type return_type_w)1728 emit_decl_sampler_view(struct ureg_program *ureg,
1729 unsigned index,
1730 enum tgsi_texture_type target,
1731 enum tgsi_return_type return_type_x,
1732 enum tgsi_return_type return_type_y,
1733 enum tgsi_return_type return_type_z,
1734 enum tgsi_return_type return_type_w )
1735 {
1736 union tgsi_any_token *out = get_tokens(ureg, DOMAIN_DECL, 3);
1737
1738 out[0].value = 0;
1739 out[0].decl.Type = TGSI_TOKEN_TYPE_DECLARATION;
1740 out[0].decl.NrTokens = 3;
1741 out[0].decl.File = TGSI_FILE_SAMPLER_VIEW;
1742 out[0].decl.UsageMask = TGSI_WRITEMASK_XYZW;
1743
1744 out[1].value = 0;
1745 out[1].decl_range.First = index;
1746 out[1].decl_range.Last = index;
1747
1748 out[2].value = 0;
1749 out[2].decl_sampler_view.Resource = target;
1750 out[2].decl_sampler_view.ReturnTypeX = return_type_x;
1751 out[2].decl_sampler_view.ReturnTypeY = return_type_y;
1752 out[2].decl_sampler_view.ReturnTypeZ = return_type_z;
1753 out[2].decl_sampler_view.ReturnTypeW = return_type_w;
1754 }
1755
1756 static void
emit_decl_image(struct ureg_program * ureg,unsigned index,enum tgsi_texture_type target,enum pipe_format format,bool wr,bool raw)1757 emit_decl_image(struct ureg_program *ureg,
1758 unsigned index,
1759 enum tgsi_texture_type target,
1760 enum pipe_format format,
1761 bool wr,
1762 bool raw)
1763 {
1764 union tgsi_any_token *out = get_tokens(ureg, DOMAIN_DECL, 3);
1765
1766 out[0].value = 0;
1767 out[0].decl.Type = TGSI_TOKEN_TYPE_DECLARATION;
1768 out[0].decl.NrTokens = 3;
1769 out[0].decl.File = TGSI_FILE_IMAGE;
1770 out[0].decl.UsageMask = TGSI_WRITEMASK_XYZW;
1771
1772 out[1].value = 0;
1773 out[1].decl_range.First = index;
1774 out[1].decl_range.Last = index;
1775
1776 out[2].value = 0;
1777 out[2].decl_image.Resource = target;
1778 out[2].decl_image.Writable = wr;
1779 out[2].decl_image.Raw = raw;
1780 out[2].decl_image.Format = format;
1781 }
1782
1783 static void
emit_decl_buffer(struct ureg_program * ureg,unsigned index,bool atomic)1784 emit_decl_buffer(struct ureg_program *ureg,
1785 unsigned index,
1786 bool atomic)
1787 {
1788 union tgsi_any_token *out = get_tokens(ureg, DOMAIN_DECL, 2);
1789
1790 out[0].value = 0;
1791 out[0].decl.Type = TGSI_TOKEN_TYPE_DECLARATION;
1792 out[0].decl.NrTokens = 2;
1793 out[0].decl.File = TGSI_FILE_BUFFER;
1794 out[0].decl.UsageMask = TGSI_WRITEMASK_XYZW;
1795 out[0].decl.Atomic = atomic;
1796
1797 out[1].value = 0;
1798 out[1].decl_range.First = index;
1799 out[1].decl_range.Last = index;
1800 }
1801
1802 static void
emit_decl_memory(struct ureg_program * ureg,unsigned memory_type)1803 emit_decl_memory(struct ureg_program *ureg, unsigned memory_type)
1804 {
1805 union tgsi_any_token *out = get_tokens(ureg, DOMAIN_DECL, 2);
1806
1807 out[0].value = 0;
1808 out[0].decl.Type = TGSI_TOKEN_TYPE_DECLARATION;
1809 out[0].decl.NrTokens = 2;
1810 out[0].decl.File = TGSI_FILE_MEMORY;
1811 out[0].decl.UsageMask = TGSI_WRITEMASK_XYZW;
1812 out[0].decl.MemType = memory_type;
1813
1814 out[1].value = 0;
1815 out[1].decl_range.First = memory_type;
1816 out[1].decl_range.Last = memory_type;
1817 }
1818
1819 static void
emit_immediate(struct ureg_program * ureg,const unsigned * v,unsigned type)1820 emit_immediate( struct ureg_program *ureg,
1821 const unsigned *v,
1822 unsigned type )
1823 {
1824 union tgsi_any_token *out = get_tokens( ureg, DOMAIN_DECL, 5 );
1825
1826 out[0].value = 0;
1827 out[0].imm.Type = TGSI_TOKEN_TYPE_IMMEDIATE;
1828 out[0].imm.NrTokens = 5;
1829 out[0].imm.DataType = type;
1830 out[0].imm.Padding = 0;
1831
1832 out[1].imm_data.Uint = v[0];
1833 out[2].imm_data.Uint = v[1];
1834 out[3].imm_data.Uint = v[2];
1835 out[4].imm_data.Uint = v[3];
1836 }
1837
1838 static void
emit_property(struct ureg_program * ureg,unsigned name,unsigned data)1839 emit_property(struct ureg_program *ureg,
1840 unsigned name,
1841 unsigned data)
1842 {
1843 union tgsi_any_token *out = get_tokens(ureg, DOMAIN_DECL, 2);
1844
1845 out[0].value = 0;
1846 out[0].prop.Type = TGSI_TOKEN_TYPE_PROPERTY;
1847 out[0].prop.NrTokens = 2;
1848 out[0].prop.PropertyName = name;
1849
1850 out[1].prop_data.Data = data;
1851 }
1852
1853 static int
input_sort(const void * in_a,const void * in_b)1854 input_sort(const void *in_a, const void *in_b)
1855 {
1856 const struct ureg_input_decl *a = in_a, *b = in_b;
1857
1858 return a->first - b->first;
1859 }
1860
1861 static int
output_sort(const void * in_a,const void * in_b)1862 output_sort(const void *in_a, const void *in_b)
1863 {
1864 const struct ureg_output_decl *a = in_a, *b = in_b;
1865
1866 return a->first - b->first;
1867 }
1868
1869 static int
atomic_decl_range_sort(const void * in_a,const void * in_b)1870 atomic_decl_range_sort(const void *in_a, const void *in_b)
1871 {
1872 const struct hw_atomic_decl_range *a = in_a, *b = in_b;
1873
1874 return a->first - b->first;
1875 }
1876
emit_decls(struct ureg_program * ureg)1877 static void emit_decls( struct ureg_program *ureg )
1878 {
1879 unsigned i,j;
1880
1881 for (i = 0; i < ARRAY_SIZE(ureg->properties); i++)
1882 if (ureg->properties[i] != ~0u)
1883 emit_property(ureg, i, ureg->properties[i]);
1884
1885 /* While not required by TGSI spec, virglrenderer has a dependency on the
1886 * inputs being sorted.
1887 */
1888 qsort(ureg->input, ureg->nr_inputs, sizeof(ureg->input[0]), input_sort);
1889
1890 if (ureg->processor == PIPE_SHADER_VERTEX) {
1891 for (i = 0; i < PIPE_MAX_ATTRIBS; i++) {
1892 if (ureg->vs_inputs[i/32] & (1u << (i%32))) {
1893 emit_decl_range( ureg, TGSI_FILE_INPUT, i, 1 );
1894 }
1895 }
1896 } else if (ureg->processor == PIPE_SHADER_FRAGMENT) {
1897 if (ureg->supports_any_inout_decl_range) {
1898 for (i = 0; i < ureg->nr_inputs; i++) {
1899 emit_decl_fs(ureg,
1900 TGSI_FILE_INPUT,
1901 ureg->input[i].first,
1902 ureg->input[i].last,
1903 ureg->input[i].semantic_name,
1904 ureg->input[i].semantic_index,
1905 ureg->input[i].interp,
1906 ureg->input[i].interp_location,
1907 ureg->input[i].array_id,
1908 ureg->input[i].usage_mask);
1909 }
1910 }
1911 else {
1912 for (i = 0; i < ureg->nr_inputs; i++) {
1913 for (j = ureg->input[i].first; j <= ureg->input[i].last; j++) {
1914 emit_decl_fs(ureg,
1915 TGSI_FILE_INPUT,
1916 j, j,
1917 ureg->input[i].semantic_name,
1918 ureg->input[i].semantic_index +
1919 (j - ureg->input[i].first),
1920 ureg->input[i].interp,
1921 ureg->input[i].interp_location, 0,
1922 ureg->input[i].usage_mask);
1923 }
1924 }
1925 }
1926 } else {
1927 if (ureg->supports_any_inout_decl_range) {
1928 for (i = 0; i < ureg->nr_inputs; i++) {
1929 emit_decl_semantic(ureg,
1930 TGSI_FILE_INPUT,
1931 ureg->input[i].first,
1932 ureg->input[i].last,
1933 ureg->input[i].semantic_name,
1934 ureg->input[i].semantic_index,
1935 0,
1936 TGSI_WRITEMASK_XYZW,
1937 ureg->input[i].array_id,
1938 false);
1939 }
1940 }
1941 else {
1942 for (i = 0; i < ureg->nr_inputs; i++) {
1943 for (j = ureg->input[i].first; j <= ureg->input[i].last; j++) {
1944 emit_decl_semantic(ureg,
1945 TGSI_FILE_INPUT,
1946 j, j,
1947 ureg->input[i].semantic_name,
1948 ureg->input[i].semantic_index +
1949 (j - ureg->input[i].first),
1950 0,
1951 TGSI_WRITEMASK_XYZW, 0, false);
1952 }
1953 }
1954 }
1955 }
1956
1957 for (i = 0; i < ureg->nr_system_values; i++) {
1958 emit_decl_semantic(ureg,
1959 TGSI_FILE_SYSTEM_VALUE,
1960 i,
1961 i,
1962 ureg->system_value[i].semantic_name,
1963 ureg->system_value[i].semantic_index,
1964 0,
1965 TGSI_WRITEMASK_XYZW, 0, false);
1966 }
1967
1968 /* While not required by TGSI spec, virglrenderer has a dependency on the
1969 * outputs being sorted.
1970 */
1971 qsort(ureg->output, ureg->nr_outputs, sizeof(ureg->output[0]), output_sort);
1972
1973 if (ureg->supports_any_inout_decl_range) {
1974 for (i = 0; i < ureg->nr_outputs; i++) {
1975 emit_decl_semantic_typed(ureg,
1976 TGSI_FILE_OUTPUT,
1977 ureg->output[i].first,
1978 ureg->output[i].last,
1979 ureg->output[i].semantic_name,
1980 ureg->output[i].semantic_index,
1981 ureg->output[i].streams,
1982 ureg->output[i].usage_mask,
1983 ureg->output[i].array_id,
1984 ureg->output[i].invariant,
1985 ureg->output[i].value_type);
1986
1987 }
1988 }
1989 else {
1990 for (i = 0; i < ureg->nr_outputs; i++) {
1991 for (j = ureg->output[i].first; j <= ureg->output[i].last; j++) {
1992 emit_decl_semantic_typed(ureg,
1993 TGSI_FILE_OUTPUT,
1994 j, j,
1995 ureg->output[i].semantic_name,
1996 ureg->output[i].semantic_index +
1997 (j - ureg->output[i].first),
1998 ureg->output[i].streams,
1999 ureg->output[i].usage_mask,
2000 0,
2001 ureg->output[i].invariant,
2002 ureg->output[i].value_type);
2003 }
2004 }
2005 }
2006
2007 for (i = 0; i < ureg->nr_samplers; i++) {
2008 emit_decl_range( ureg,
2009 TGSI_FILE_SAMPLER,
2010 ureg->sampler[i].Index, 1 );
2011 }
2012
2013 for (i = 0; i < ureg->nr_sampler_views; i++) {
2014 emit_decl_sampler_view(ureg,
2015 ureg->sampler_view[i].index,
2016 ureg->sampler_view[i].target,
2017 ureg->sampler_view[i].return_type_x,
2018 ureg->sampler_view[i].return_type_y,
2019 ureg->sampler_view[i].return_type_z,
2020 ureg->sampler_view[i].return_type_w);
2021 }
2022
2023 for (i = 0; i < ureg->nr_images; i++) {
2024 emit_decl_image(ureg,
2025 ureg->image[i].index,
2026 ureg->image[i].target,
2027 ureg->image[i].format,
2028 ureg->image[i].wr,
2029 ureg->image[i].raw);
2030 }
2031
2032 for (i = 0; i < ureg->nr_buffers; i++) {
2033 emit_decl_buffer(ureg, ureg->buffer[i].index, ureg->buffer[i].atomic);
2034 }
2035
2036 for (i = 0; i < TGSI_MEMORY_TYPE_COUNT; i++) {
2037 if (ureg->use_memory[i])
2038 emit_decl_memory(ureg, i);
2039 }
2040
2041 for (i = 0; i < PIPE_MAX_CONSTANT_BUFFERS; i++) {
2042 struct const_decl *decl = &ureg->const_decls[i];
2043
2044 if (decl->nr_constant_ranges) {
2045 unsigned j;
2046
2047 for (j = 0; j < decl->nr_constant_ranges; j++) {
2048 emit_decl_range2D(ureg,
2049 TGSI_FILE_CONSTANT,
2050 decl->constant_range[j].first,
2051 decl->constant_range[j].last,
2052 i);
2053 }
2054 }
2055 }
2056
2057 for (i = 0; i < PIPE_MAX_HW_ATOMIC_BUFFERS; i++) {
2058 struct hw_atomic_decl *decl = &ureg->hw_atomic_decls[i];
2059
2060 if (decl->nr_hw_atomic_ranges) {
2061 unsigned j;
2062
2063 /* GLSL-to-TGSI generated HW atomic counters in order, and r600 depends
2064 * on it.
2065 */
2066 qsort(decl->hw_atomic_range, decl->nr_hw_atomic_ranges, sizeof(struct hw_atomic_decl_range), atomic_decl_range_sort);
2067
2068 for (j = 0; j < decl->nr_hw_atomic_ranges; j++) {
2069 emit_decl_atomic_2d(ureg,
2070 decl->hw_atomic_range[j].first,
2071 decl->hw_atomic_range[j].last,
2072 i,
2073 decl->hw_atomic_range[j].array_id);
2074 }
2075 }
2076 }
2077
2078 if (ureg->nr_temps) {
2079 unsigned array = 0;
2080 for (i = 0; i < ureg->nr_temps;) {
2081 bool local = util_bitmask_get(ureg->local_temps, i);
2082 unsigned first = i;
2083 i = util_bitmask_get_next_index(ureg->decl_temps, i + 1);
2084 if (i == UTIL_BITMASK_INVALID_INDEX)
2085 i = ureg->nr_temps;
2086
2087 if (array < ureg->nr_array_temps && ureg->array_temps[array] == first)
2088 emit_decl_temps( ureg, first, i - 1, local, ++array );
2089 else
2090 emit_decl_temps( ureg, first, i - 1, local, 0 );
2091 }
2092 }
2093
2094 if (ureg->nr_addrs) {
2095 emit_decl_range( ureg,
2096 TGSI_FILE_ADDRESS,
2097 0, ureg->nr_addrs );
2098 }
2099
2100 for (i = 0; i < ureg->nr_immediates; i++) {
2101 emit_immediate( ureg,
2102 ureg->immediate[i].value.u,
2103 ureg->immediate[i].type );
2104 }
2105 }
2106
2107 /* Append the instruction tokens onto the declarations to build a
2108 * contiguous stream suitable to send to the driver.
2109 */
copy_instructions(struct ureg_program * ureg)2110 static void copy_instructions( struct ureg_program *ureg )
2111 {
2112 unsigned nr_tokens = ureg->domain[DOMAIN_INSN].count;
2113 union tgsi_any_token *out = get_tokens( ureg,
2114 DOMAIN_DECL,
2115 nr_tokens );
2116
2117 memcpy(out,
2118 ureg->domain[DOMAIN_INSN].tokens,
2119 nr_tokens * sizeof out[0] );
2120 }
2121
2122
2123 static void
fixup_header_size(struct ureg_program * ureg)2124 fixup_header_size(struct ureg_program *ureg)
2125 {
2126 union tgsi_any_token *out = retrieve_token( ureg, DOMAIN_DECL, 0 );
2127
2128 out->header.BodySize = ureg->domain[DOMAIN_DECL].count - 2;
2129 }
2130
2131
2132 static void
emit_header(struct ureg_program * ureg)2133 emit_header( struct ureg_program *ureg )
2134 {
2135 union tgsi_any_token *out = get_tokens( ureg, DOMAIN_DECL, 2 );
2136
2137 out[0].header.HeaderSize = 2;
2138 out[0].header.BodySize = 0;
2139
2140 out[1].processor.Processor = ureg->processor;
2141 out[1].processor.Padding = 0;
2142 }
2143
2144
ureg_finalize(struct ureg_program * ureg)2145 const struct tgsi_token *ureg_finalize( struct ureg_program *ureg )
2146 {
2147 const struct tgsi_token *tokens;
2148
2149 switch (ureg->processor) {
2150 case PIPE_SHADER_VERTEX:
2151 case PIPE_SHADER_TESS_EVAL:
2152 ureg_property(ureg, TGSI_PROPERTY_NEXT_SHADER,
2153 ureg->next_shader_processor == -1 ?
2154 PIPE_SHADER_FRAGMENT :
2155 ureg->next_shader_processor);
2156 break;
2157 default:
2158 ; /* nothing */
2159 }
2160
2161 emit_header( ureg );
2162 emit_decls( ureg );
2163 copy_instructions( ureg );
2164 fixup_header_size( ureg );
2165
2166 if (ureg->domain[0].tokens == error_tokens ||
2167 ureg->domain[1].tokens == error_tokens) {
2168 debug_printf("%s: error in generated shader\n", __func__);
2169 assert(0);
2170 return NULL;
2171 }
2172
2173 tokens = &ureg->domain[DOMAIN_DECL].tokens[0].token;
2174
2175 if (0) {
2176 debug_printf("%s: emitted shader %d tokens:\n", __func__,
2177 ureg->domain[DOMAIN_DECL].count);
2178 tgsi_dump( tokens, 0 );
2179 }
2180
2181 #if MESA_DEBUG
2182 /* tgsi_sanity doesn't seem to return if there are too many constants. */
2183 bool too_many_constants = false;
2184 for (unsigned i = 0; i < ARRAY_SIZE(ureg->const_decls); i++) {
2185 for (unsigned j = 0; j < ureg->const_decls[i].nr_constant_ranges; j++) {
2186 if (ureg->const_decls[i].constant_range[j].last > 4096) {
2187 too_many_constants = true;
2188 break;
2189 }
2190 }
2191 }
2192
2193 if (tokens && !too_many_constants && !tgsi_sanity_check(tokens)) {
2194 debug_printf("tgsi_ureg.c, sanity check failed on generated tokens:\n");
2195 tgsi_dump(tokens, 0);
2196 assert(0);
2197 }
2198 #endif
2199
2200
2201 return tokens;
2202 }
2203
2204
ureg_create_shader(struct ureg_program * ureg,struct pipe_context * pipe,const struct pipe_stream_output_info * so)2205 void *ureg_create_shader( struct ureg_program *ureg,
2206 struct pipe_context *pipe,
2207 const struct pipe_stream_output_info *so )
2208 {
2209 struct pipe_shader_state state = {0};
2210
2211 pipe_shader_state_from_tgsi(&state, ureg_finalize(ureg));
2212 if(!state.tokens)
2213 return NULL;
2214
2215 if (so)
2216 state.stream_output = *so;
2217
2218 switch (ureg->processor) {
2219 case PIPE_SHADER_VERTEX:
2220 return pipe->create_vs_state(pipe, &state);
2221 case PIPE_SHADER_TESS_CTRL:
2222 return pipe->create_tcs_state(pipe, &state);
2223 case PIPE_SHADER_TESS_EVAL:
2224 return pipe->create_tes_state(pipe, &state);
2225 case PIPE_SHADER_GEOMETRY:
2226 return pipe->create_gs_state(pipe, &state);
2227 case PIPE_SHADER_FRAGMENT:
2228 return pipe->create_fs_state(pipe, &state);
2229 default:
2230 return NULL;
2231 }
2232 }
2233
2234
ureg_get_tokens(struct ureg_program * ureg,unsigned * nr_tokens)2235 const struct tgsi_token *ureg_get_tokens( struct ureg_program *ureg,
2236 unsigned *nr_tokens )
2237 {
2238 const struct tgsi_token *tokens;
2239
2240 ureg_finalize(ureg);
2241
2242 tokens = &ureg->domain[DOMAIN_DECL].tokens[0].token;
2243
2244 if (nr_tokens)
2245 *nr_tokens = ureg->domain[DOMAIN_DECL].count;
2246
2247 ureg->domain[DOMAIN_DECL].tokens = NULL;
2248 ureg->domain[DOMAIN_DECL].size = 0;
2249 ureg->domain[DOMAIN_DECL].order = 0;
2250 ureg->domain[DOMAIN_DECL].count = 0;
2251
2252 return tokens;
2253 }
2254
2255
ureg_free_tokens(const struct tgsi_token * tokens)2256 void ureg_free_tokens( const struct tgsi_token *tokens )
2257 {
2258 FREE((struct tgsi_token *)tokens);
2259 }
2260
2261
2262 struct ureg_program *
ureg_create(enum pipe_shader_type processor)2263 ureg_create(enum pipe_shader_type processor)
2264 {
2265 return ureg_create_with_screen(processor, NULL);
2266 }
2267
2268
2269 struct ureg_program *
ureg_create_with_screen(enum pipe_shader_type processor,struct pipe_screen * screen)2270 ureg_create_with_screen(enum pipe_shader_type processor,
2271 struct pipe_screen *screen)
2272 {
2273 unsigned i;
2274 struct ureg_program *ureg = CALLOC_STRUCT( ureg_program );
2275 if (!ureg)
2276 goto no_ureg;
2277
2278 ureg->processor = processor;
2279 ureg->supports_any_inout_decl_range =
2280 screen &&
2281 screen->get_shader_param(screen, processor,
2282 PIPE_SHADER_CAP_TGSI_ANY_INOUT_DECL_RANGE) != 0;
2283 ureg->next_shader_processor = -1;
2284
2285 for (i = 0; i < ARRAY_SIZE(ureg->properties); i++)
2286 ureg->properties[i] = ~0;
2287
2288 ureg->free_temps = util_bitmask_create();
2289 if (ureg->free_temps == NULL)
2290 goto no_free_temps;
2291
2292 ureg->local_temps = util_bitmask_create();
2293 if (ureg->local_temps == NULL)
2294 goto no_local_temps;
2295
2296 ureg->decl_temps = util_bitmask_create();
2297 if (ureg->decl_temps == NULL)
2298 goto no_decl_temps;
2299
2300 return ureg;
2301
2302 no_decl_temps:
2303 util_bitmask_destroy(ureg->local_temps);
2304 no_local_temps:
2305 util_bitmask_destroy(ureg->free_temps);
2306 no_free_temps:
2307 FREE(ureg);
2308 no_ureg:
2309 return NULL;
2310 }
2311
2312
2313 void
ureg_set_next_shader_processor(struct ureg_program * ureg,unsigned processor)2314 ureg_set_next_shader_processor(struct ureg_program *ureg, unsigned processor)
2315 {
2316 ureg->next_shader_processor = processor;
2317 }
2318
2319
2320 unsigned
ureg_get_nr_outputs(const struct ureg_program * ureg)2321 ureg_get_nr_outputs( const struct ureg_program *ureg )
2322 {
2323 if (!ureg)
2324 return 0;
2325 return ureg->nr_outputs;
2326 }
2327
2328 static void
ureg_setup_clipdist_info(struct ureg_program * ureg,const struct shader_info * info)2329 ureg_setup_clipdist_info(struct ureg_program *ureg,
2330 const struct shader_info *info)
2331 {
2332 if (info->clip_distance_array_size)
2333 ureg_property(ureg, TGSI_PROPERTY_NUM_CLIPDIST_ENABLED,
2334 info->clip_distance_array_size);
2335 if (info->cull_distance_array_size)
2336 ureg_property(ureg, TGSI_PROPERTY_NUM_CULLDIST_ENABLED,
2337 info->cull_distance_array_size);
2338 }
2339
2340 static void
ureg_setup_tess_ctrl_shader(struct ureg_program * ureg,const struct shader_info * info)2341 ureg_setup_tess_ctrl_shader(struct ureg_program *ureg,
2342 const struct shader_info *info)
2343 {
2344 ureg_property(ureg, TGSI_PROPERTY_TCS_VERTICES_OUT,
2345 info->tess.tcs_vertices_out);
2346 }
2347
2348 static void
ureg_setup_tess_eval_shader(struct ureg_program * ureg,const struct shader_info * info)2349 ureg_setup_tess_eval_shader(struct ureg_program *ureg,
2350 const struct shader_info *info)
2351 {
2352 ureg_property(ureg, TGSI_PROPERTY_TES_PRIM_MODE, u_tess_prim_from_shader(info->tess._primitive_mode));
2353
2354 STATIC_ASSERT((TESS_SPACING_EQUAL + 1) % 3 == PIPE_TESS_SPACING_EQUAL);
2355 STATIC_ASSERT((TESS_SPACING_FRACTIONAL_ODD + 1) % 3 ==
2356 PIPE_TESS_SPACING_FRACTIONAL_ODD);
2357 STATIC_ASSERT((TESS_SPACING_FRACTIONAL_EVEN + 1) % 3 ==
2358 PIPE_TESS_SPACING_FRACTIONAL_EVEN);
2359
2360 ureg_property(ureg, TGSI_PROPERTY_TES_SPACING,
2361 (info->tess.spacing + 1) % 3);
2362
2363 ureg_property(ureg, TGSI_PROPERTY_TES_VERTEX_ORDER_CW,
2364 !info->tess.ccw);
2365 ureg_property(ureg, TGSI_PROPERTY_TES_POINT_MODE,
2366 info->tess.point_mode);
2367 }
2368
2369 static void
ureg_setup_geometry_shader(struct ureg_program * ureg,const struct shader_info * info)2370 ureg_setup_geometry_shader(struct ureg_program *ureg,
2371 const struct shader_info *info)
2372 {
2373 ureg_property(ureg, TGSI_PROPERTY_GS_INPUT_PRIM,
2374 info->gs.input_primitive);
2375 ureg_property(ureg, TGSI_PROPERTY_GS_OUTPUT_PRIM,
2376 info->gs.output_primitive);
2377 ureg_property(ureg, TGSI_PROPERTY_GS_MAX_OUTPUT_VERTICES,
2378 info->gs.vertices_out);
2379 ureg_property(ureg, TGSI_PROPERTY_GS_INVOCATIONS,
2380 info->gs.invocations);
2381 }
2382
2383 static void
ureg_setup_fragment_shader(struct ureg_program * ureg,const struct shader_info * info)2384 ureg_setup_fragment_shader(struct ureg_program *ureg,
2385 const struct shader_info *info)
2386 {
2387 if (info->fs.early_fragment_tests || info->fs.post_depth_coverage) {
2388 ureg_property(ureg, TGSI_PROPERTY_FS_EARLY_DEPTH_STENCIL, 1);
2389
2390 if (info->fs.post_depth_coverage)
2391 ureg_property(ureg, TGSI_PROPERTY_FS_POST_DEPTH_COVERAGE, 1);
2392 }
2393
2394 if (info->fs.depth_layout != FRAG_DEPTH_LAYOUT_NONE) {
2395 switch (info->fs.depth_layout) {
2396 case FRAG_DEPTH_LAYOUT_ANY:
2397 ureg_property(ureg, TGSI_PROPERTY_FS_DEPTH_LAYOUT,
2398 TGSI_FS_DEPTH_LAYOUT_ANY);
2399 break;
2400 case FRAG_DEPTH_LAYOUT_GREATER:
2401 ureg_property(ureg, TGSI_PROPERTY_FS_DEPTH_LAYOUT,
2402 TGSI_FS_DEPTH_LAYOUT_GREATER);
2403 break;
2404 case FRAG_DEPTH_LAYOUT_LESS:
2405 ureg_property(ureg, TGSI_PROPERTY_FS_DEPTH_LAYOUT,
2406 TGSI_FS_DEPTH_LAYOUT_LESS);
2407 break;
2408 case FRAG_DEPTH_LAYOUT_UNCHANGED:
2409 ureg_property(ureg, TGSI_PROPERTY_FS_DEPTH_LAYOUT,
2410 TGSI_FS_DEPTH_LAYOUT_UNCHANGED);
2411 break;
2412 default:
2413 assert(0);
2414 }
2415 }
2416
2417 if (info->fs.advanced_blend_modes) {
2418 ureg_property(ureg, TGSI_PROPERTY_FS_BLEND_EQUATION_ADVANCED,
2419 info->fs.advanced_blend_modes);
2420 }
2421 }
2422
2423 static void
ureg_setup_compute_shader(struct ureg_program * ureg,const struct shader_info * info)2424 ureg_setup_compute_shader(struct ureg_program *ureg,
2425 const struct shader_info *info)
2426 {
2427 ureg_property(ureg, TGSI_PROPERTY_CS_FIXED_BLOCK_WIDTH,
2428 info->workgroup_size[0]);
2429 ureg_property(ureg, TGSI_PROPERTY_CS_FIXED_BLOCK_HEIGHT,
2430 info->workgroup_size[1]);
2431 ureg_property(ureg, TGSI_PROPERTY_CS_FIXED_BLOCK_DEPTH,
2432 info->workgroup_size[2]);
2433
2434 if (info->shared_size)
2435 ureg_DECL_memory(ureg, TGSI_MEMORY_TYPE_SHARED);
2436 }
2437
2438 void
ureg_setup_shader_info(struct ureg_program * ureg,const struct shader_info * info)2439 ureg_setup_shader_info(struct ureg_program *ureg,
2440 const struct shader_info *info)
2441 {
2442 if (info->layer_viewport_relative)
2443 ureg_property(ureg, TGSI_PROPERTY_LAYER_VIEWPORT_RELATIVE, 1);
2444
2445 switch (info->stage) {
2446 case MESA_SHADER_VERTEX:
2447 ureg_setup_clipdist_info(ureg, info);
2448 ureg_set_next_shader_processor(ureg, pipe_shader_type_from_mesa(info->next_stage));
2449 break;
2450 case MESA_SHADER_TESS_CTRL:
2451 ureg_setup_tess_ctrl_shader(ureg, info);
2452 break;
2453 case MESA_SHADER_TESS_EVAL:
2454 ureg_setup_tess_eval_shader(ureg, info);
2455 ureg_setup_clipdist_info(ureg, info);
2456 ureg_set_next_shader_processor(ureg, pipe_shader_type_from_mesa(info->next_stage));
2457 break;
2458 case MESA_SHADER_GEOMETRY:
2459 ureg_setup_geometry_shader(ureg, info);
2460 ureg_setup_clipdist_info(ureg, info);
2461 break;
2462 case MESA_SHADER_FRAGMENT:
2463 ureg_setup_fragment_shader(ureg, info);
2464 break;
2465 case MESA_SHADER_COMPUTE:
2466 ureg_setup_compute_shader(ureg, info);
2467 break;
2468 default:
2469 break;
2470 }
2471 }
2472
2473
ureg_destroy(struct ureg_program * ureg)2474 void ureg_destroy( struct ureg_program *ureg )
2475 {
2476 unsigned i;
2477
2478 for (i = 0; i < ARRAY_SIZE(ureg->domain); i++) {
2479 if (ureg->domain[i].tokens &&
2480 ureg->domain[i].tokens != error_tokens)
2481 FREE(ureg->domain[i].tokens);
2482 }
2483
2484 util_bitmask_destroy(ureg->free_temps);
2485 util_bitmask_destroy(ureg->local_temps);
2486 util_bitmask_destroy(ureg->decl_temps);
2487
2488 FREE(ureg);
2489 }
2490
ureg_set_precise(struct ureg_program * ureg,bool precise)2491 void ureg_set_precise( struct ureg_program *ureg, bool precise )
2492 {
2493 ureg->precise = precise;
2494 }
2495