1 /*
2 * Copyright © 2014 Connor Abbott
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 DEALINGS
21 * IN THE SOFTWARE.
22 *
23 * Authors:
24 * Connor Abbott (cwabbott0@gmail.com)
25 *
26 */
27
28 #ifndef NIR_H
29 #define NIR_H
30
31 #include "util/hash_table.h"
32 #include "compiler/glsl/list.h"
33 #include "util/list.h"
34 #include "util/log.h"
35 #include "util/ralloc.h"
36 #include "util/set.h"
37 #include "util/bitscan.h"
38 #include "util/bitset.h"
39 #include "util/compiler.h"
40 #include "util/enum_operators.h"
41 #include "util/macros.h"
42 #include "util/format/u_format.h"
43 #include "compiler/nir_types.h"
44 #include "compiler/shader_enums.h"
45 #include "compiler/shader_info.h"
46 #define XXH_INLINE_ALL
47 #include "util/xxhash.h"
48 #include <stdio.h>
49
50 #ifndef NDEBUG
51 #include "util/debug.h"
52 #endif /* NDEBUG */
53
54 #include "nir_opcodes.h"
55
56 #if defined(_WIN32) && !defined(snprintf)
57 #define snprintf _snprintf
58 #endif
59
60 #ifdef __cplusplus
61 extern "C" {
62 #endif
63
64 extern uint32_t nir_debug;
65 extern bool nir_debug_print_shader[MESA_SHADER_KERNEL + 1];
66
67 #ifndef NDEBUG
68 #define NIR_DEBUG(flag) unlikely(nir_debug & (NIR_DEBUG_ ## flag))
69 #else
70 #define NIR_DEBUG(flag) false
71 #endif
72
73 #define NIR_DEBUG_CLONE (1u << 0)
74 #define NIR_DEBUG_SERIALIZE (1u << 1)
75 #define NIR_DEBUG_NOVALIDATE (1u << 2)
76 #define NIR_DEBUG_VALIDATE_SSA_DOMINANCE (1u << 3)
77 #define NIR_DEBUG_TGSI (1u << 4)
78 #define NIR_DEBUG_PRINT_VS (1u << 5)
79 #define NIR_DEBUG_PRINT_TCS (1u << 6)
80 #define NIR_DEBUG_PRINT_TES (1u << 7)
81 #define NIR_DEBUG_PRINT_GS (1u << 8)
82 #define NIR_DEBUG_PRINT_FS (1u << 9)
83 #define NIR_DEBUG_PRINT_CS (1u << 10)
84 #define NIR_DEBUG_PRINT_TS (1u << 11)
85 #define NIR_DEBUG_PRINT_MS (1u << 12)
86 #define NIR_DEBUG_PRINT_RGS (1u << 13)
87 #define NIR_DEBUG_PRINT_AHS (1u << 14)
88 #define NIR_DEBUG_PRINT_CHS (1u << 15)
89 #define NIR_DEBUG_PRINT_MHS (1u << 16)
90 #define NIR_DEBUG_PRINT_IS (1u << 17)
91 #define NIR_DEBUG_PRINT_CBS (1u << 18)
92 #define NIR_DEBUG_PRINT_KS (1u << 19)
93 #define NIR_DEBUG_PRINT_CONSTS (1u << 20)
94 #define NIR_DEBUG_VALIDATE_GC_LIST (1u << 21)
95
96 #define NIR_DEBUG_PRINT (NIR_DEBUG_PRINT_VS | \
97 NIR_DEBUG_PRINT_TCS | \
98 NIR_DEBUG_PRINT_TES | \
99 NIR_DEBUG_PRINT_GS | \
100 NIR_DEBUG_PRINT_FS | \
101 NIR_DEBUG_PRINT_CS | \
102 NIR_DEBUG_PRINT_TS | \
103 NIR_DEBUG_PRINT_MS | \
104 NIR_DEBUG_PRINT_RGS | \
105 NIR_DEBUG_PRINT_AHS | \
106 NIR_DEBUG_PRINT_CHS | \
107 NIR_DEBUG_PRINT_MHS | \
108 NIR_DEBUG_PRINT_IS | \
109 NIR_DEBUG_PRINT_CBS | \
110 NIR_DEBUG_PRINT_KS)
111
112 #define NIR_FALSE 0u
113 #define NIR_TRUE (~0u)
114 #define NIR_MAX_VEC_COMPONENTS 16
115 #define NIR_MAX_MATRIX_COLUMNS 4
116 #define NIR_STREAM_PACKED (1 << 8)
117 typedef uint16_t nir_component_mask_t;
118
119 static inline bool
nir_num_components_valid(unsigned num_components)120 nir_num_components_valid(unsigned num_components)
121 {
122 return (num_components >= 1 &&
123 num_components <= 5) ||
124 num_components == 8 ||
125 num_components == 16;
126 }
127
128 void
129 nir_process_debug_variable(void);
130
131 bool nir_component_mask_can_reinterpret(nir_component_mask_t mask,
132 unsigned old_bit_size,
133 unsigned new_bit_size);
134 nir_component_mask_t
135 nir_component_mask_reinterpret(nir_component_mask_t mask,
136 unsigned old_bit_size,
137 unsigned new_bit_size);
138
139 /** Defines a cast function
140 *
141 * This macro defines a cast function from in_type to out_type where
142 * out_type is some structure type that contains a field of type out_type.
143 *
144 * Note that you have to be a bit careful as the generated cast function
145 * destroys constness.
146 */
147 #define NIR_DEFINE_CAST(name, in_type, out_type, field, \
148 type_field, type_value) \
149 static inline out_type * \
150 name(const in_type *parent) \
151 { \
152 assert(parent && parent->type_field == type_value); \
153 return exec_node_data(out_type, parent, field); \
154 }
155
156 struct nir_function;
157 struct nir_shader;
158 struct nir_instr;
159 struct nir_builder;
160 struct nir_xfb_info;
161
162
163 /**
164 * Description of built-in state associated with a uniform
165 *
166 * \sa nir_variable::state_slots
167 */
168 typedef struct {
169 gl_state_index16 tokens[STATE_LENGTH];
170 uint16_t swizzle;
171 } nir_state_slot;
172
173 typedef enum {
174 nir_var_system_value = (1 << 0),
175 nir_var_uniform = (1 << 1),
176 nir_var_shader_in = (1 << 2),
177 nir_var_shader_out = (1 << 3),
178 nir_var_image = (1 << 4),
179 /** Incoming call or ray payload data for ray-tracing shaders */
180 nir_var_shader_call_data = (1 << 5),
181 /** Ray hit attributes */
182 nir_var_ray_hit_attrib = (1 << 6),
183
184 /* Modes named nir_var_mem_* have explicit data layout */
185 nir_var_mem_ubo = (1 << 7),
186 nir_var_mem_push_const = (1 << 8),
187 nir_var_mem_ssbo = (1 << 9),
188 nir_var_mem_constant = (1 << 10),
189 nir_var_mem_task_payload = (1 << 11),
190
191 /* Generic modes intentionally come last. See encode_dref_modes() in
192 * nir_serialize.c for more details.
193 */
194 nir_var_shader_temp = (1 << 12),
195 nir_var_function_temp = (1 << 13),
196 nir_var_mem_shared = (1 << 14),
197 nir_var_mem_global = (1 << 15),
198
199 nir_var_mem_generic = (nir_var_shader_temp |
200 nir_var_function_temp |
201 nir_var_mem_shared |
202 nir_var_mem_global),
203
204 nir_var_read_only_modes = nir_var_shader_in | nir_var_uniform |
205 nir_var_system_value | nir_var_mem_constant |
206 nir_var_mem_ubo,
207 /** Modes where vector derefs can be indexed as arrays */
208 nir_var_vec_indexable_modes = nir_var_mem_ubo | nir_var_mem_ssbo |
209 nir_var_mem_shared | nir_var_mem_global |
210 nir_var_mem_push_const,
211 nir_num_variable_modes = 16,
212 nir_var_all = (1 << nir_num_variable_modes) - 1,
213 } nir_variable_mode;
214 MESA_DEFINE_CPP_ENUM_BITFIELD_OPERATORS(nir_variable_mode)
215
216 /**
217 * Rounding modes.
218 */
219 typedef enum {
220 nir_rounding_mode_undef = 0,
221 nir_rounding_mode_rtne = 1, /* round to nearest even */
222 nir_rounding_mode_ru = 2, /* round up */
223 nir_rounding_mode_rd = 3, /* round down */
224 nir_rounding_mode_rtz = 4, /* round towards zero */
225 } nir_rounding_mode;
226
227 typedef union {
228 bool b;
229 float f32;
230 double f64;
231 int8_t i8;
232 uint8_t u8;
233 int16_t i16;
234 uint16_t u16;
235 int32_t i32;
236 uint32_t u32;
237 int64_t i64;
238 uint64_t u64;
239 } nir_const_value;
240
241 #define nir_const_value_to_array(arr, c, components, m) \
242 { \
243 for (unsigned i = 0; i < components; ++i) \
244 arr[i] = c[i].m; \
245 } while (false)
246
247 static inline nir_const_value
nir_const_value_for_raw_uint(uint64_t x,unsigned bit_size)248 nir_const_value_for_raw_uint(uint64_t x, unsigned bit_size)
249 {
250 nir_const_value v;
251 memset(&v, 0, sizeof(v));
252
253 switch (bit_size) {
254 case 1: v.b = x; break;
255 case 8: v.u8 = x; break;
256 case 16: v.u16 = x; break;
257 case 32: v.u32 = x; break;
258 case 64: v.u64 = x; break;
259 default:
260 unreachable("Invalid bit size");
261 }
262
263 return v;
264 }
265
266 static inline nir_const_value
nir_const_value_for_int(int64_t i,unsigned bit_size)267 nir_const_value_for_int(int64_t i, unsigned bit_size)
268 {
269 nir_const_value v;
270 memset(&v, 0, sizeof(v));
271
272 assert(bit_size <= 64);
273 if (bit_size < 64) {
274 assert(i >= (-(1ll << (bit_size - 1))));
275 assert(i < (1ll << (bit_size - 1)));
276 }
277
278 return nir_const_value_for_raw_uint(i, bit_size);
279 }
280
281 static inline nir_const_value
nir_const_value_for_uint(uint64_t u,unsigned bit_size)282 nir_const_value_for_uint(uint64_t u, unsigned bit_size)
283 {
284 nir_const_value v;
285 memset(&v, 0, sizeof(v));
286
287 assert(bit_size <= 64);
288 if (bit_size < 64)
289 assert(u < (1ull << bit_size));
290
291 return nir_const_value_for_raw_uint(u, bit_size);
292 }
293
294 static inline nir_const_value
nir_const_value_for_bool(bool b,unsigned bit_size)295 nir_const_value_for_bool(bool b, unsigned bit_size)
296 {
297 /* Booleans use a 0/-1 convention */
298 return nir_const_value_for_int(-(int)b, bit_size);
299 }
300
301 /* This one isn't inline because it requires half-float conversion */
302 nir_const_value nir_const_value_for_float(double b, unsigned bit_size);
303
304 static inline int64_t
nir_const_value_as_int(nir_const_value value,unsigned bit_size)305 nir_const_value_as_int(nir_const_value value, unsigned bit_size)
306 {
307 switch (bit_size) {
308 /* int1_t uses 0/-1 convention */
309 case 1: return -(int)value.b;
310 case 8: return value.i8;
311 case 16: return value.i16;
312 case 32: return value.i32;
313 case 64: return value.i64;
314 default:
315 unreachable("Invalid bit size");
316 }
317 }
318
319 static inline uint64_t
nir_const_value_as_uint(nir_const_value value,unsigned bit_size)320 nir_const_value_as_uint(nir_const_value value, unsigned bit_size)
321 {
322 switch (bit_size) {
323 case 1: return value.b;
324 case 8: return value.u8;
325 case 16: return value.u16;
326 case 32: return value.u32;
327 case 64: return value.u64;
328 default:
329 unreachable("Invalid bit size");
330 }
331 }
332
333 static inline bool
nir_const_value_as_bool(nir_const_value value,unsigned bit_size)334 nir_const_value_as_bool(nir_const_value value, unsigned bit_size)
335 {
336 int64_t i = nir_const_value_as_int(value, bit_size);
337
338 /* Booleans of any size use 0/-1 convention */
339 assert(i == 0 || i == -1);
340
341 return i;
342 }
343
344 /* This one isn't inline because it requires half-float conversion */
345 double nir_const_value_as_float(nir_const_value value, unsigned bit_size);
346
347 typedef struct nir_constant {
348 /**
349 * Value of the constant.
350 *
351 * The field used to back the values supplied by the constant is determined
352 * by the type associated with the \c nir_variable. Constants may be
353 * scalars, vectors, or matrices.
354 */
355 nir_const_value values[NIR_MAX_VEC_COMPONENTS];
356
357 /* we could get this from the var->type but makes clone *much* easier to
358 * not have to care about the type.
359 */
360 unsigned num_elements;
361
362 /* Array elements / Structure Fields */
363 struct nir_constant **elements;
364 } nir_constant;
365
366 /**
367 * \brief Layout qualifiers for gl_FragDepth.
368 *
369 * The AMD/ARB_conservative_depth extensions allow gl_FragDepth to be redeclared
370 * with a layout qualifier.
371 */
372 typedef enum {
373 nir_depth_layout_none, /**< No depth layout is specified. */
374 nir_depth_layout_any,
375 nir_depth_layout_greater,
376 nir_depth_layout_less,
377 nir_depth_layout_unchanged
378 } nir_depth_layout;
379
380 /**
381 * Enum keeping track of how a variable was declared.
382 */
383 typedef enum {
384 /**
385 * Normal declaration.
386 */
387 nir_var_declared_normally = 0,
388
389 /**
390 * Variable is implicitly generated by the compiler and should not be
391 * visible via the API.
392 */
393 nir_var_hidden,
394 } nir_var_declaration_type;
395
396 /**
397 * Either a uniform, global variable, shader input, or shader output. Based on
398 * ir_variable - it should be easy to translate between the two.
399 */
400
401 typedef struct nir_variable {
402 struct exec_node node;
403
404 /**
405 * Declared type of the variable
406 */
407 const struct glsl_type *type;
408
409 /**
410 * Declared name of the variable
411 */
412 char *name;
413
414 struct nir_variable_data {
415 /**
416 * Storage class of the variable.
417 *
418 * \sa nir_variable_mode
419 */
420 unsigned mode:15;
421
422 /**
423 * Is the variable read-only?
424 *
425 * This is set for variables declared as \c const, shader inputs,
426 * and uniforms.
427 */
428 unsigned read_only:1;
429 unsigned centroid:1;
430 unsigned sample:1;
431 unsigned patch:1;
432 unsigned invariant:1;
433
434 /**
435 * Is the variable a ray query?
436 */
437 unsigned ray_query:1;
438
439 /**
440 * Precision qualifier.
441 *
442 * In desktop GLSL we do not care about precision qualifiers at all, in
443 * fact, the spec says that precision qualifiers are ignored.
444 *
445 * To make things easy, we make it so that this field is always
446 * GLSL_PRECISION_NONE on desktop shaders. This way all the variables
447 * have the same precision value and the checks we add in the compiler
448 * for this field will never break a desktop shader compile.
449 */
450 unsigned precision:2;
451
452 /**
453 * Can this variable be coalesced with another?
454 *
455 * This is set by nir_lower_io_to_temporaries to say that any
456 * copies involving this variable should stay put. Propagating it can
457 * duplicate the resulting load/store, which is not wanted, and may
458 * result in a load/store of the variable with an indirect offset which
459 * the backend may not be able to handle.
460 */
461 unsigned cannot_coalesce:1;
462
463 /**
464 * When separate shader programs are enabled, only input/outputs between
465 * the stages of a multi-stage separate program can be safely removed
466 * from the shader interface. Other input/outputs must remains active.
467 *
468 * This is also used to make sure xfb varyings that are unused by the
469 * fragment shader are not removed.
470 */
471 unsigned always_active_io:1;
472
473 /**
474 * Interpolation mode for shader inputs / outputs
475 *
476 * \sa glsl_interp_mode
477 */
478 unsigned interpolation:3;
479
480 /**
481 * If non-zero, then this variable may be packed along with other variables
482 * into a single varying slot, so this offset should be applied when
483 * accessing components. For example, an offset of 1 means that the x
484 * component of this variable is actually stored in component y of the
485 * location specified by \c location.
486 */
487 unsigned location_frac:2;
488
489 /**
490 * If true, this variable represents an array of scalars that should
491 * be tightly packed. In other words, consecutive array elements
492 * should be stored one component apart, rather than one slot apart.
493 */
494 unsigned compact:1;
495
496 /**
497 * Whether this is a fragment shader output implicitly initialized with
498 * the previous contents of the specified render target at the
499 * framebuffer location corresponding to this shader invocation.
500 */
501 unsigned fb_fetch_output:1;
502
503 /**
504 * Non-zero if this variable is considered bindless as defined by
505 * ARB_bindless_texture.
506 */
507 unsigned bindless:1;
508
509 /**
510 * Was an explicit binding set in the shader?
511 */
512 unsigned explicit_binding:1;
513
514 /**
515 * Was the location explicitly set in the shader?
516 *
517 * If the location is explicitly set in the shader, it \b cannot be changed
518 * by the linker or by the API (e.g., calls to \c glBindAttribLocation have
519 * no effect).
520 */
521 unsigned explicit_location:1;
522
523 /**
524 * Was a transfer feedback buffer set in the shader?
525 */
526 unsigned explicit_xfb_buffer:1;
527
528 /**
529 * Was a transfer feedback stride set in the shader?
530 */
531 unsigned explicit_xfb_stride:1;
532
533 /**
534 * Was an explicit offset set in the shader?
535 */
536 unsigned explicit_offset:1;
537
538 /**
539 * Layout of the matrix. Uses glsl_matrix_layout values.
540 */
541 unsigned matrix_layout:2;
542
543 /**
544 * Non-zero if this variable was created by lowering a named interface
545 * block.
546 */
547 unsigned from_named_ifc_block:1;
548
549 /**
550 * How the variable was declared. See nir_var_declaration_type.
551 *
552 * This is used to detect variables generated by the compiler, so should
553 * not be visible via the API.
554 */
555 unsigned how_declared:2;
556
557 /**
558 * Is this variable per-view? If so, we know it must be an array with
559 * size corresponding to the number of views.
560 */
561 unsigned per_view:1;
562
563 /**
564 * Whether the variable is per-primitive.
565 * Can be use by Mesh Shader outputs and corresponding Fragment Shader inputs.
566 */
567 unsigned per_primitive:1;
568
569 /**
570 * \brief Layout qualifier for gl_FragDepth. See nir_depth_layout.
571 *
572 * This is not equal to \c ir_depth_layout_none if and only if this
573 * variable is \c gl_FragDepth and a layout qualifier is specified.
574 */
575 unsigned depth_layout:3;
576
577 /**
578 * Vertex stream output identifier.
579 *
580 * For packed outputs, NIR_STREAM_PACKED is set and bits [2*i+1,2*i]
581 * indicate the stream of the i-th component.
582 */
583 unsigned stream:9;
584
585 /**
586 * See gl_access_qualifier.
587 *
588 * Access flags for memory variables (SSBO/global), image uniforms, and
589 * bindless images in uniforms/inputs/outputs.
590 */
591 unsigned access:9;
592
593 /**
594 * Descriptor set binding for sampler or UBO.
595 */
596 unsigned descriptor_set:5;
597
598 /**
599 * output index for dual source blending.
600 */
601 unsigned index;
602
603 /**
604 * Initial binding point for a sampler or UBO.
605 *
606 * For array types, this represents the binding point for the first element.
607 */
608 unsigned binding;
609
610 /**
611 * Storage location of the base of this variable
612 *
613 * The precise meaning of this field depends on the nature of the variable.
614 *
615 * - Vertex shader input: one of the values from \c gl_vert_attrib.
616 * - Vertex shader output: one of the values from \c gl_varying_slot.
617 * - Geometry shader input: one of the values from \c gl_varying_slot.
618 * - Geometry shader output: one of the values from \c gl_varying_slot.
619 * - Fragment shader input: one of the values from \c gl_varying_slot.
620 * - Fragment shader output: one of the values from \c gl_frag_result.
621 * - Task shader output: one of the values from \c gl_varying_slot.
622 * - Mesh shader input: one of the values from \c gl_varying_slot.
623 * - Mesh shader output: one of the values from \c gl_varying_slot.
624 * - Uniforms: Per-stage uniform slot number for default uniform block.
625 * - Uniforms: Index within the uniform block definition for UBO members.
626 * - Non-UBO Uniforms: uniform slot number.
627 * - Other: This field is not currently used.
628 *
629 * If the variable is a uniform, shader input, or shader output, and the
630 * slot has not been assigned, the value will be -1.
631 */
632 int location;
633
634 /**
635 * The actual location of the variable in the IR. Only valid for inputs,
636 * outputs, uniforms (including samplers and images), and for UBO and SSBO
637 * variables in GLSL.
638 */
639 unsigned driver_location;
640
641 /**
642 * Location an atomic counter or transform feedback is stored at.
643 */
644 unsigned offset;
645
646 union {
647 struct {
648 /** Image internal format if specified explicitly, otherwise PIPE_FORMAT_NONE. */
649 enum pipe_format format;
650 } image;
651
652 struct {
653 /**
654 * For OpenCL inline samplers. See cl_sampler_addressing_mode and cl_sampler_filter_mode
655 */
656 unsigned is_inline_sampler : 1;
657 unsigned addressing_mode : 3;
658 unsigned normalized_coordinates : 1;
659 unsigned filter_mode : 1;
660 } sampler;
661
662 struct {
663 /**
664 * Transform feedback buffer.
665 */
666 uint16_t buffer:2;
667
668 /**
669 * Transform feedback stride.
670 */
671 uint16_t stride;
672 } xfb;
673 };
674 } data;
675
676 /**
677 * Identifier for this variable generated by nir_index_vars() that is unique
678 * among other variables in the same exec_list.
679 */
680 unsigned index;
681
682 /* Number of nir_variable_data members */
683 uint16_t num_members;
684
685 /**
686 * Built-in state that backs this uniform
687 *
688 * Once set at variable creation, \c state_slots must remain invariant.
689 * This is because, ideally, this array would be shared by all clones of
690 * this variable in the IR tree. In other words, we'd really like for it
691 * to be a fly-weight.
692 *
693 * If the variable is not a uniform, \c num_state_slots will be zero and
694 * \c state_slots will be \c NULL.
695 */
696 /*@{*/
697 uint16_t num_state_slots; /**< Number of state slots used */
698 nir_state_slot *state_slots; /**< State descriptors. */
699 /*@}*/
700
701 /**
702 * Constant expression assigned in the initializer of the variable
703 *
704 * This field should only be used temporarily by creators of NIR shaders
705 * and then nir_lower_variable_initializers can be used to get rid of them.
706 * Most of the rest of NIR ignores this field or asserts that it's NULL.
707 */
708 nir_constant *constant_initializer;
709
710 /**
711 * Global variable assigned in the initializer of the variable
712 * This field should only be used temporarily by creators of NIR shaders
713 * and then nir_lower_variable_initializers can be used to get rid of them.
714 * Most of the rest of NIR ignores this field or asserts that it's NULL.
715 */
716 struct nir_variable *pointer_initializer;
717
718 /**
719 * For variables that are in an interface block or are an instance of an
720 * interface block, this is the \c GLSL_TYPE_INTERFACE type for that block.
721 *
722 * \sa ir_variable::location
723 */
724 const struct glsl_type *interface_type;
725
726 /**
727 * Description of per-member data for per-member struct variables
728 *
729 * This is used for variables which are actually an amalgamation of
730 * multiple entities such as a struct of built-in values or a struct of
731 * inputs each with their own layout specifier. This is only allowed on
732 * variables with a struct or array of array of struct type.
733 */
734 struct nir_variable_data *members;
735 } nir_variable;
736
737 static inline bool
_nir_shader_variable_has_mode(nir_variable * var,unsigned modes)738 _nir_shader_variable_has_mode(nir_variable *var, unsigned modes)
739 {
740 /* This isn't a shader variable */
741 assert(!(modes & nir_var_function_temp));
742 return var->data.mode & modes;
743 }
744
745 #define nir_foreach_variable_in_list(var, var_list) \
746 foreach_list_typed(nir_variable, var, node, var_list)
747
748 #define nir_foreach_variable_in_list_safe(var, var_list) \
749 foreach_list_typed_safe(nir_variable, var, node, var_list)
750
751 #define nir_foreach_variable_in_shader(var, shader) \
752 nir_foreach_variable_in_list(var, &(shader)->variables)
753
754 #define nir_foreach_variable_in_shader_safe(var, shader) \
755 nir_foreach_variable_in_list_safe(var, &(shader)->variables)
756
757 #define nir_foreach_variable_with_modes(var, shader, modes) \
758 nir_foreach_variable_in_shader(var, shader) \
759 if (_nir_shader_variable_has_mode(var, modes))
760
761 #define nir_foreach_variable_with_modes_safe(var, shader, modes) \
762 nir_foreach_variable_in_shader_safe(var, shader) \
763 if (_nir_shader_variable_has_mode(var, modes))
764
765 #define nir_foreach_shader_in_variable(var, shader) \
766 nir_foreach_variable_with_modes(var, shader, nir_var_shader_in)
767
768 #define nir_foreach_shader_in_variable_safe(var, shader) \
769 nir_foreach_variable_with_modes_safe(var, shader, nir_var_shader_in)
770
771 #define nir_foreach_shader_out_variable(var, shader) \
772 nir_foreach_variable_with_modes(var, shader, nir_var_shader_out)
773
774 #define nir_foreach_shader_out_variable_safe(var, shader) \
775 nir_foreach_variable_with_modes_safe(var, shader, nir_var_shader_out)
776
777 #define nir_foreach_uniform_variable(var, shader) \
778 nir_foreach_variable_with_modes(var, shader, nir_var_uniform)
779
780 #define nir_foreach_uniform_variable_safe(var, shader) \
781 nir_foreach_variable_with_modes_safe(var, shader, nir_var_uniform)
782
783 #define nir_foreach_image_variable(var, shader) \
784 nir_foreach_variable_with_modes(var, shader, nir_var_image)
785
786 #define nir_foreach_image_variable_safe(var, shader) \
787 nir_foreach_variable_with_modes_safe(var, shader, nir_var_image)
788
789 static inline bool
nir_variable_is_global(const nir_variable * var)790 nir_variable_is_global(const nir_variable *var)
791 {
792 return var->data.mode != nir_var_function_temp;
793 }
794
795 typedef struct nir_register {
796 struct exec_node node;
797
798 unsigned num_components; /** < number of vector components */
799 unsigned num_array_elems; /** < size of array (0 for no array) */
800
801 /* The bit-size of each channel; must be one of 8, 16, 32, or 64 */
802 uint8_t bit_size;
803
804 /**
805 * True if this register may have different values in different SIMD
806 * invocations of the shader.
807 */
808 bool divergent;
809
810 /** generic register index. */
811 unsigned index;
812
813 /** set of nir_srcs where this register is used (read from) */
814 struct list_head uses;
815
816 /** set of nir_dests where this register is defined (written to) */
817 struct list_head defs;
818
819 /** set of nir_ifs where this register is used as a condition */
820 struct list_head if_uses;
821 } nir_register;
822
823 #define nir_foreach_register(reg, reg_list) \
824 foreach_list_typed(nir_register, reg, node, reg_list)
825 #define nir_foreach_register_safe(reg, reg_list) \
826 foreach_list_typed_safe(nir_register, reg, node, reg_list)
827
828 typedef enum PACKED {
829 nir_instr_type_alu,
830 nir_instr_type_deref,
831 nir_instr_type_call,
832 nir_instr_type_tex,
833 nir_instr_type_intrinsic,
834 nir_instr_type_load_const,
835 nir_instr_type_jump,
836 nir_instr_type_ssa_undef,
837 nir_instr_type_phi,
838 nir_instr_type_parallel_copy,
839 } nir_instr_type;
840
841 typedef struct nir_instr {
842 struct exec_node node;
843 struct list_head gc_node;
844 struct nir_block *block;
845 nir_instr_type type;
846
847 /* A temporary for optimization and analysis passes to use for storing
848 * flags. For instance, DCE uses this to store the "dead/live" info.
849 */
850 uint8_t pass_flags;
851
852 /** generic instruction index. */
853 uint32_t index;
854 } nir_instr;
855
856 static inline nir_instr *
nir_instr_next(nir_instr * instr)857 nir_instr_next(nir_instr *instr)
858 {
859 struct exec_node *next = exec_node_get_next(&instr->node);
860 if (exec_node_is_tail_sentinel(next))
861 return NULL;
862 else
863 return exec_node_data(nir_instr, next, node);
864 }
865
866 static inline nir_instr *
nir_instr_prev(nir_instr * instr)867 nir_instr_prev(nir_instr *instr)
868 {
869 struct exec_node *prev = exec_node_get_prev(&instr->node);
870 if (exec_node_is_head_sentinel(prev))
871 return NULL;
872 else
873 return exec_node_data(nir_instr, prev, node);
874 }
875
876 static inline bool
nir_instr_is_first(const nir_instr * instr)877 nir_instr_is_first(const nir_instr *instr)
878 {
879 return exec_node_is_head_sentinel(exec_node_get_prev_const(&instr->node));
880 }
881
882 static inline bool
nir_instr_is_last(const nir_instr * instr)883 nir_instr_is_last(const nir_instr *instr)
884 {
885 return exec_node_is_tail_sentinel(exec_node_get_next_const(&instr->node));
886 }
887
888 typedef struct nir_ssa_def {
889 /** Instruction which produces this SSA value. */
890 nir_instr *parent_instr;
891
892 /** set of nir_instrs where this register is used (read from) */
893 struct list_head uses;
894
895 /** set of nir_ifs where this register is used as a condition */
896 struct list_head if_uses;
897
898 /** generic SSA definition index. */
899 unsigned index;
900
901 uint8_t num_components;
902
903 /* The bit-size of each channel; must be one of 8, 16, 32, or 64 */
904 uint8_t bit_size;
905
906 /**
907 * True if this SSA value may have different values in different SIMD
908 * invocations of the shader. This is set by nir_divergence_analysis.
909 */
910 bool divergent;
911 } nir_ssa_def;
912
913 struct nir_src;
914
915 typedef struct {
916 nir_register *reg;
917 struct nir_src *indirect; /** < NULL for no indirect offset */
918 unsigned base_offset;
919
920 /* TODO use-def chain goes here */
921 } nir_reg_src;
922
923 typedef struct {
924 nir_instr *parent_instr;
925 struct list_head def_link;
926
927 nir_register *reg;
928 struct nir_src *indirect; /** < NULL for no indirect offset */
929 unsigned base_offset;
930
931 /* TODO def-use chain goes here */
932 } nir_reg_dest;
933
934 struct nir_if;
935
936 typedef struct nir_src {
937 union {
938 /** Instruction that consumes this value as a source. */
939 nir_instr *parent_instr;
940 struct nir_if *parent_if;
941 };
942
943 struct list_head use_link;
944
945 union {
946 nir_reg_src reg;
947 nir_ssa_def *ssa;
948 };
949
950 bool is_ssa;
951 } nir_src;
952
953 static inline nir_src
nir_src_init(void)954 nir_src_init(void)
955 {
956 nir_src src = { { NULL } };
957 return src;
958 }
959
960 #define NIR_SRC_INIT nir_src_init()
961
962 #define nir_foreach_use(src, reg_or_ssa_def) \
963 list_for_each_entry(nir_src, src, &(reg_or_ssa_def)->uses, use_link)
964
965 #define nir_foreach_use_safe(src, reg_or_ssa_def) \
966 list_for_each_entry_safe(nir_src, src, &(reg_or_ssa_def)->uses, use_link)
967
968 #define nir_foreach_if_use(src, reg_or_ssa_def) \
969 list_for_each_entry(nir_src, src, &(reg_or_ssa_def)->if_uses, use_link)
970
971 #define nir_foreach_if_use_safe(src, reg_or_ssa_def) \
972 list_for_each_entry_safe(nir_src, src, &(reg_or_ssa_def)->if_uses, use_link)
973
974 typedef struct {
975 union {
976 nir_reg_dest reg;
977 nir_ssa_def ssa;
978 };
979
980 bool is_ssa;
981 } nir_dest;
982
983 static inline nir_dest
nir_dest_init(void)984 nir_dest_init(void)
985 {
986 nir_dest dest = { { { NULL } } };
987 return dest;
988 }
989
990 #define NIR_DEST_INIT nir_dest_init()
991
992 #define nir_foreach_def(dest, reg) \
993 list_for_each_entry(nir_dest, dest, &(reg)->defs, reg.def_link)
994
995 #define nir_foreach_def_safe(dest, reg) \
996 list_for_each_entry_safe(nir_dest, dest, &(reg)->defs, reg.def_link)
997
998 static inline nir_src
nir_src_for_ssa(nir_ssa_def * def)999 nir_src_for_ssa(nir_ssa_def *def)
1000 {
1001 nir_src src = NIR_SRC_INIT;
1002
1003 src.is_ssa = true;
1004 src.ssa = def;
1005
1006 return src;
1007 }
1008
1009 static inline nir_src
nir_src_for_reg(nir_register * reg)1010 nir_src_for_reg(nir_register *reg)
1011 {
1012 nir_src src = NIR_SRC_INIT;
1013
1014 src.is_ssa = false;
1015 src.reg.reg = reg;
1016 src.reg.indirect = NULL;
1017 src.reg.base_offset = 0;
1018
1019 return src;
1020 }
1021
1022 static inline nir_dest
nir_dest_for_reg(nir_register * reg)1023 nir_dest_for_reg(nir_register *reg)
1024 {
1025 nir_dest dest = NIR_DEST_INIT;
1026
1027 dest.reg.reg = reg;
1028
1029 return dest;
1030 }
1031
1032 static inline unsigned
nir_src_bit_size(nir_src src)1033 nir_src_bit_size(nir_src src)
1034 {
1035 return src.is_ssa ? src.ssa->bit_size : src.reg.reg->bit_size;
1036 }
1037
1038 static inline unsigned
nir_src_num_components(nir_src src)1039 nir_src_num_components(nir_src src)
1040 {
1041 return src.is_ssa ? src.ssa->num_components : src.reg.reg->num_components;
1042 }
1043
1044 static inline bool
nir_src_is_const(nir_src src)1045 nir_src_is_const(nir_src src)
1046 {
1047 return src.is_ssa &&
1048 src.ssa->parent_instr->type == nir_instr_type_load_const;
1049 }
1050
1051 static inline bool
nir_src_is_undef(nir_src src)1052 nir_src_is_undef(nir_src src)
1053 {
1054 return src.is_ssa &&
1055 src.ssa->parent_instr->type == nir_instr_type_ssa_undef;
1056 }
1057
1058 static inline bool
nir_src_is_divergent(nir_src src)1059 nir_src_is_divergent(nir_src src)
1060 {
1061 return src.is_ssa ? src.ssa->divergent : src.reg.reg->divergent;
1062 }
1063
1064 static inline unsigned
nir_dest_bit_size(nir_dest dest)1065 nir_dest_bit_size(nir_dest dest)
1066 {
1067 return dest.is_ssa ? dest.ssa.bit_size : dest.reg.reg->bit_size;
1068 }
1069
1070 static inline unsigned
nir_dest_num_components(nir_dest dest)1071 nir_dest_num_components(nir_dest dest)
1072 {
1073 return dest.is_ssa ? dest.ssa.num_components : dest.reg.reg->num_components;
1074 }
1075
1076 static inline bool
nir_dest_is_divergent(nir_dest dest)1077 nir_dest_is_divergent(nir_dest dest)
1078 {
1079 return dest.is_ssa ? dest.ssa.divergent : dest.reg.reg->divergent;
1080 }
1081
1082 /* Are all components the same, ie. .xxxx */
1083 static inline bool
nir_is_same_comp_swizzle(uint8_t * swiz,unsigned nr_comp)1084 nir_is_same_comp_swizzle(uint8_t *swiz, unsigned nr_comp)
1085 {
1086 for (unsigned i = 1; i < nr_comp; i++)
1087 if (swiz[i] != swiz[0])
1088 return false;
1089 return true;
1090 }
1091
1092 /* Are all components sequential, ie. .yzw */
1093 static inline bool
nir_is_sequential_comp_swizzle(uint8_t * swiz,unsigned nr_comp)1094 nir_is_sequential_comp_swizzle(uint8_t *swiz, unsigned nr_comp)
1095 {
1096 for (unsigned i = 1; i < nr_comp; i++)
1097 if (swiz[i] != (swiz[0] + i))
1098 return false;
1099 return true;
1100 }
1101
1102 void nir_src_copy(nir_src *dest, const nir_src *src);
1103 void nir_dest_copy(nir_dest *dest, const nir_dest *src);
1104
1105 typedef struct {
1106 /** Base source */
1107 nir_src src;
1108
1109 /**
1110 * \name input modifiers
1111 */
1112 /*@{*/
1113 /**
1114 * For inputs interpreted as floating point, flips the sign bit. For
1115 * inputs interpreted as integers, performs the two's complement negation.
1116 */
1117 bool negate;
1118
1119 /**
1120 * Clears the sign bit for floating point values, and computes the integer
1121 * absolute value for integers. Note that the negate modifier acts after
1122 * the absolute value modifier, therefore if both are set then all inputs
1123 * will become negative.
1124 */
1125 bool abs;
1126 /*@}*/
1127
1128 /**
1129 * For each input component, says which component of the register it is
1130 * chosen from.
1131 *
1132 * Note that which elements of the swizzle are used and which are ignored
1133 * are based on the write mask for most opcodes - for example, a statement
1134 * like "foo.xzw = bar.zyx" would have a writemask of 1101b and a swizzle
1135 * of {2, 1, x, 0} where x means "don't care."
1136 */
1137 uint8_t swizzle[NIR_MAX_VEC_COMPONENTS];
1138 } nir_alu_src;
1139
1140 typedef struct {
1141 /** Base destination */
1142 nir_dest dest;
1143
1144 /**
1145 * Saturate output modifier
1146 *
1147 * Only valid for opcodes that output floating-point numbers. Clamps the
1148 * output to between 0.0 and 1.0 inclusive.
1149 */
1150 bool saturate;
1151
1152 /**
1153 * Write-mask
1154 *
1155 * Ignored if dest.is_ssa is true
1156 */
1157 unsigned write_mask : NIR_MAX_VEC_COMPONENTS;
1158 } nir_alu_dest;
1159
1160 /** NIR sized and unsized types
1161 *
1162 * The values in this enum are carefully chosen so that the sized type is
1163 * just the unsized type OR the number of bits.
1164 */
1165 typedef enum PACKED {
1166 nir_type_invalid = 0, /* Not a valid type */
1167 nir_type_int = 2,
1168 nir_type_uint = 4,
1169 nir_type_bool = 6,
1170 nir_type_float = 128,
1171 nir_type_bool1 = 1 | nir_type_bool,
1172 nir_type_bool8 = 8 | nir_type_bool,
1173 nir_type_bool16 = 16 | nir_type_bool,
1174 nir_type_bool32 = 32 | nir_type_bool,
1175 nir_type_int1 = 1 | nir_type_int,
1176 nir_type_int8 = 8 | nir_type_int,
1177 nir_type_int16 = 16 | nir_type_int,
1178 nir_type_int32 = 32 | nir_type_int,
1179 nir_type_int64 = 64 | nir_type_int,
1180 nir_type_uint1 = 1 | nir_type_uint,
1181 nir_type_uint8 = 8 | nir_type_uint,
1182 nir_type_uint16 = 16 | nir_type_uint,
1183 nir_type_uint32 = 32 | nir_type_uint,
1184 nir_type_uint64 = 64 | nir_type_uint,
1185 nir_type_float16 = 16 | nir_type_float,
1186 nir_type_float32 = 32 | nir_type_float,
1187 nir_type_float64 = 64 | nir_type_float,
1188 } nir_alu_type;
1189
1190 #define NIR_ALU_TYPE_SIZE_MASK 0x79
1191 #define NIR_ALU_TYPE_BASE_TYPE_MASK 0x86
1192
1193 static inline unsigned
nir_alu_type_get_type_size(nir_alu_type type)1194 nir_alu_type_get_type_size(nir_alu_type type)
1195 {
1196 return type & NIR_ALU_TYPE_SIZE_MASK;
1197 }
1198
1199 static inline nir_alu_type
nir_alu_type_get_base_type(nir_alu_type type)1200 nir_alu_type_get_base_type(nir_alu_type type)
1201 {
1202 return (nir_alu_type)(type & NIR_ALU_TYPE_BASE_TYPE_MASK);
1203 }
1204
1205 nir_alu_type
1206 nir_get_nir_type_for_glsl_base_type(enum glsl_base_type base_type);
1207
1208 static inline nir_alu_type
nir_get_nir_type_for_glsl_type(const struct glsl_type * type)1209 nir_get_nir_type_for_glsl_type(const struct glsl_type *type)
1210 {
1211 return nir_get_nir_type_for_glsl_base_type(glsl_get_base_type(type));
1212 }
1213
1214 enum glsl_base_type
1215 nir_get_glsl_base_type_for_nir_type(nir_alu_type base_type);
1216
1217 nir_op nir_type_conversion_op(nir_alu_type src, nir_alu_type dst,
1218 nir_rounding_mode rnd);
1219
1220 nir_op
1221 nir_op_vec(unsigned components);
1222
1223 bool
1224 nir_op_is_vec(nir_op op);
1225
1226 static inline bool
nir_is_float_control_signed_zero_inf_nan_preserve(unsigned execution_mode,unsigned bit_size)1227 nir_is_float_control_signed_zero_inf_nan_preserve(unsigned execution_mode, unsigned bit_size)
1228 {
1229 return (16 == bit_size && execution_mode & FLOAT_CONTROLS_SIGNED_ZERO_INF_NAN_PRESERVE_FP16) ||
1230 (32 == bit_size && execution_mode & FLOAT_CONTROLS_SIGNED_ZERO_INF_NAN_PRESERVE_FP32) ||
1231 (64 == bit_size && execution_mode & FLOAT_CONTROLS_SIGNED_ZERO_INF_NAN_PRESERVE_FP64);
1232 }
1233
1234 static inline bool
nir_is_denorm_flush_to_zero(unsigned execution_mode,unsigned bit_size)1235 nir_is_denorm_flush_to_zero(unsigned execution_mode, unsigned bit_size)
1236 {
1237 return (16 == bit_size && execution_mode & FLOAT_CONTROLS_DENORM_FLUSH_TO_ZERO_FP16) ||
1238 (32 == bit_size && execution_mode & FLOAT_CONTROLS_DENORM_FLUSH_TO_ZERO_FP32) ||
1239 (64 == bit_size && execution_mode & FLOAT_CONTROLS_DENORM_FLUSH_TO_ZERO_FP64);
1240 }
1241
1242 static inline bool
nir_is_denorm_preserve(unsigned execution_mode,unsigned bit_size)1243 nir_is_denorm_preserve(unsigned execution_mode, unsigned bit_size)
1244 {
1245 return (16 == bit_size && execution_mode & FLOAT_CONTROLS_DENORM_PRESERVE_FP16) ||
1246 (32 == bit_size && execution_mode & FLOAT_CONTROLS_DENORM_PRESERVE_FP32) ||
1247 (64 == bit_size && execution_mode & FLOAT_CONTROLS_DENORM_PRESERVE_FP64);
1248 }
1249
1250 static inline bool
nir_is_rounding_mode_rtne(unsigned execution_mode,unsigned bit_size)1251 nir_is_rounding_mode_rtne(unsigned execution_mode, unsigned bit_size)
1252 {
1253 return (16 == bit_size && execution_mode & FLOAT_CONTROLS_ROUNDING_MODE_RTE_FP16) ||
1254 (32 == bit_size && execution_mode & FLOAT_CONTROLS_ROUNDING_MODE_RTE_FP32) ||
1255 (64 == bit_size && execution_mode & FLOAT_CONTROLS_ROUNDING_MODE_RTE_FP64);
1256 }
1257
1258 static inline bool
nir_is_rounding_mode_rtz(unsigned execution_mode,unsigned bit_size)1259 nir_is_rounding_mode_rtz(unsigned execution_mode, unsigned bit_size)
1260 {
1261 return (16 == bit_size && execution_mode & FLOAT_CONTROLS_ROUNDING_MODE_RTZ_FP16) ||
1262 (32 == bit_size && execution_mode & FLOAT_CONTROLS_ROUNDING_MODE_RTZ_FP32) ||
1263 (64 == bit_size && execution_mode & FLOAT_CONTROLS_ROUNDING_MODE_RTZ_FP64);
1264 }
1265
1266 static inline bool
nir_has_any_rounding_mode_rtz(unsigned execution_mode)1267 nir_has_any_rounding_mode_rtz(unsigned execution_mode)
1268 {
1269 return (execution_mode & FLOAT_CONTROLS_ROUNDING_MODE_RTZ_FP16) ||
1270 (execution_mode & FLOAT_CONTROLS_ROUNDING_MODE_RTZ_FP32) ||
1271 (execution_mode & FLOAT_CONTROLS_ROUNDING_MODE_RTZ_FP64);
1272 }
1273
1274 static inline bool
nir_has_any_rounding_mode_rtne(unsigned execution_mode)1275 nir_has_any_rounding_mode_rtne(unsigned execution_mode)
1276 {
1277 return (execution_mode & FLOAT_CONTROLS_ROUNDING_MODE_RTE_FP16) ||
1278 (execution_mode & FLOAT_CONTROLS_ROUNDING_MODE_RTE_FP32) ||
1279 (execution_mode & FLOAT_CONTROLS_ROUNDING_MODE_RTE_FP64);
1280 }
1281
1282 static inline nir_rounding_mode
nir_get_rounding_mode_from_float_controls(unsigned execution_mode,nir_alu_type type)1283 nir_get_rounding_mode_from_float_controls(unsigned execution_mode,
1284 nir_alu_type type)
1285 {
1286 if (nir_alu_type_get_base_type(type) != nir_type_float)
1287 return nir_rounding_mode_undef;
1288
1289 unsigned bit_size = nir_alu_type_get_type_size(type);
1290
1291 if (nir_is_rounding_mode_rtz(execution_mode, bit_size))
1292 return nir_rounding_mode_rtz;
1293 if (nir_is_rounding_mode_rtne(execution_mode, bit_size))
1294 return nir_rounding_mode_rtne;
1295 return nir_rounding_mode_undef;
1296 }
1297
1298 static inline bool
nir_has_any_rounding_mode_enabled(unsigned execution_mode)1299 nir_has_any_rounding_mode_enabled(unsigned execution_mode)
1300 {
1301 bool result =
1302 nir_has_any_rounding_mode_rtne(execution_mode) ||
1303 nir_has_any_rounding_mode_rtz(execution_mode);
1304 return result;
1305 }
1306
1307 typedef enum {
1308 /**
1309 * Operation where the first two sources are commutative.
1310 *
1311 * For 2-source operations, this just mathematical commutativity. Some
1312 * 3-source operations, like ffma, are only commutative in the first two
1313 * sources.
1314 */
1315 NIR_OP_IS_2SRC_COMMUTATIVE = (1 << 0),
1316
1317 /**
1318 * Operation is associative
1319 */
1320 NIR_OP_IS_ASSOCIATIVE = (1 << 1),
1321 } nir_op_algebraic_property;
1322
1323 /* vec16 is the widest ALU op in NIR, making the max number of input of ALU
1324 * instructions to be the same as NIR_MAX_VEC_COMPONENTS.
1325 */
1326 #define NIR_ALU_MAX_INPUTS NIR_MAX_VEC_COMPONENTS
1327
1328 typedef struct nir_op_info {
1329 /** Name of the NIR ALU opcode */
1330 const char *name;
1331
1332 /** Number of inputs (sources) */
1333 uint8_t num_inputs;
1334
1335 /**
1336 * The number of components in the output
1337 *
1338 * If non-zero, this is the size of the output and input sizes are
1339 * explicitly given; swizzle and writemask are still in effect, but if
1340 * the output component is masked out, then the input component may
1341 * still be in use.
1342 *
1343 * If zero, the opcode acts in the standard, per-component manner; the
1344 * operation is performed on each component (except the ones that are
1345 * masked out) with the input being taken from the input swizzle for
1346 * that component.
1347 *
1348 * The size of some of the inputs may be given (i.e. non-zero) even
1349 * though output_size is zero; in that case, the inputs with a zero
1350 * size act per-component, while the inputs with non-zero size don't.
1351 */
1352 uint8_t output_size;
1353
1354 /**
1355 * The type of vector that the instruction outputs. Note that the
1356 * staurate modifier is only allowed on outputs with the float type.
1357 */
1358 nir_alu_type output_type;
1359
1360 /**
1361 * The number of components in each input
1362 *
1363 * See nir_op_infos::output_size for more detail about the relationship
1364 * between input and output sizes.
1365 */
1366 uint8_t input_sizes[NIR_ALU_MAX_INPUTS];
1367
1368 /**
1369 * The type of vector that each input takes. Note that negate and
1370 * absolute value are only allowed on inputs with int or float type and
1371 * behave differently on the two.
1372 */
1373 nir_alu_type input_types[NIR_ALU_MAX_INPUTS];
1374
1375 /** Algebraic properties of this opcode */
1376 nir_op_algebraic_property algebraic_properties;
1377
1378 /** Whether this represents a numeric conversion opcode */
1379 bool is_conversion;
1380 } nir_op_info;
1381
1382 /** Metadata for each nir_op, indexed by opcode */
1383 extern const nir_op_info nir_op_infos[nir_num_opcodes];
1384
1385 typedef struct nir_alu_instr {
1386 /** Base instruction */
1387 nir_instr instr;
1388
1389 /** Opcode */
1390 nir_op op;
1391
1392 /** Indicates that this ALU instruction generates an exact value
1393 *
1394 * This is kind of a mixture of GLSL "precise" and "invariant" and not
1395 * really equivalent to either. This indicates that the value generated by
1396 * this operation is high-precision and any code transformations that touch
1397 * it must ensure that the resulting value is bit-for-bit identical to the
1398 * original.
1399 */
1400 bool exact:1;
1401
1402 /**
1403 * Indicates that this instruction doese not cause signed integer wrapping
1404 * to occur, in the form of overflow or underflow.
1405 */
1406 bool no_signed_wrap:1;
1407
1408 /**
1409 * Indicates that this instruction does not cause unsigned integer wrapping
1410 * to occur, in the form of overflow or underflow.
1411 */
1412 bool no_unsigned_wrap:1;
1413
1414 /** Destination */
1415 nir_alu_dest dest;
1416
1417 /** Sources
1418 *
1419 * The size of the array is given by nir_op_info::num_inputs.
1420 */
1421 nir_alu_src src[];
1422 } nir_alu_instr;
1423
1424 void nir_alu_src_copy(nir_alu_src *dest, const nir_alu_src *src);
1425 void nir_alu_dest_copy(nir_alu_dest *dest, const nir_alu_dest *src);
1426
1427 bool nir_alu_instr_is_copy(nir_alu_instr *instr);
1428
1429 /* is this source channel used? */
1430 bool
1431 nir_alu_instr_channel_used(const nir_alu_instr *instr, unsigned src,
1432 unsigned channel);
1433 nir_component_mask_t
1434 nir_alu_instr_src_read_mask(const nir_alu_instr *instr, unsigned src);
1435 /**
1436 * Get the number of channels used for a source
1437 */
1438 unsigned
1439 nir_ssa_alu_instr_src_components(const nir_alu_instr *instr, unsigned src);
1440
1441 bool
1442 nir_alu_instr_is_comparison(const nir_alu_instr *instr);
1443
1444 bool nir_const_value_negative_equal(nir_const_value c1, nir_const_value c2,
1445 nir_alu_type full_type);
1446
1447 bool nir_alu_srcs_equal(const nir_alu_instr *alu1, const nir_alu_instr *alu2,
1448 unsigned src1, unsigned src2);
1449
1450 bool nir_alu_srcs_negative_equal(const nir_alu_instr *alu1,
1451 const nir_alu_instr *alu2,
1452 unsigned src1, unsigned src2);
1453
1454 bool nir_alu_src_is_trivial_ssa(const nir_alu_instr *alu, unsigned srcn);
1455
1456 typedef enum {
1457 nir_deref_type_var,
1458 nir_deref_type_array,
1459 nir_deref_type_array_wildcard,
1460 nir_deref_type_ptr_as_array,
1461 nir_deref_type_struct,
1462 nir_deref_type_cast,
1463 } nir_deref_type;
1464
1465 typedef struct {
1466 nir_instr instr;
1467
1468 /** The type of this deref instruction */
1469 nir_deref_type deref_type;
1470
1471 /** Bitmask what modes the underlying variable might be
1472 *
1473 * For OpenCL-style generic pointers, we may not know exactly what mode it
1474 * is at any given point in time in the compile process. This bitfield
1475 * contains the set of modes which it MAY be.
1476 *
1477 * Generally, this field should not be accessed directly. Use one of the
1478 * nir_deref_mode_ helpers instead.
1479 */
1480 nir_variable_mode modes;
1481
1482 /** The dereferenced type of the resulting pointer value */
1483 const struct glsl_type *type;
1484
1485 union {
1486 /** Variable being dereferenced if deref_type is a deref_var */
1487 nir_variable *var;
1488
1489 /** Parent deref if deref_type is not deref_var */
1490 nir_src parent;
1491 };
1492
1493 /** Additional deref parameters */
1494 union {
1495 struct {
1496 nir_src index;
1497 } arr;
1498
1499 struct {
1500 unsigned index;
1501 } strct;
1502
1503 struct {
1504 unsigned ptr_stride;
1505 unsigned align_mul;
1506 unsigned align_offset;
1507 } cast;
1508 };
1509
1510 /** Destination to store the resulting "pointer" */
1511 nir_dest dest;
1512 } nir_deref_instr;
1513
1514 /** Returns true if deref might have one of the given modes
1515 *
1516 * For multi-mode derefs, this returns true if any of the possible modes for
1517 * the deref to have any of the specified modes. This function returning true
1518 * does NOT mean that the deref definitely has one of those modes. It simply
1519 * means that, with the best information we have at the time, it might.
1520 */
1521 static inline bool
nir_deref_mode_may_be(const nir_deref_instr * deref,nir_variable_mode modes)1522 nir_deref_mode_may_be(const nir_deref_instr *deref, nir_variable_mode modes)
1523 {
1524 assert(!(modes & ~nir_var_all));
1525 assert(deref->modes != 0);
1526 return deref->modes & modes;
1527 }
1528
1529 /** Returns true if deref must have one of the given modes
1530 *
1531 * For multi-mode derefs, this returns true if NIR can prove that the given
1532 * deref has one of the specified modes. This function returning false does
1533 * NOT mean that deref doesn't have one of the given mode. It very well may
1534 * have one of those modes, we just don't have enough information to prove
1535 * that it does for sure.
1536 */
1537 static inline bool
nir_deref_mode_must_be(const nir_deref_instr * deref,nir_variable_mode modes)1538 nir_deref_mode_must_be(const nir_deref_instr *deref, nir_variable_mode modes)
1539 {
1540 assert(!(modes & ~nir_var_all));
1541 assert(deref->modes != 0);
1542 return !(deref->modes & ~modes);
1543 }
1544
1545 /** Returns true if deref has the given mode
1546 *
1547 * This returns true if the deref has exactly the mode specified. If the
1548 * deref may have that mode but may also have a different mode (i.e. modes has
1549 * multiple bits set), this will assert-fail.
1550 *
1551 * If you're confused about which nir_deref_mode_ helper to use, use this one
1552 * or nir_deref_mode_is_one_of below.
1553 */
1554 static inline bool
nir_deref_mode_is(const nir_deref_instr * deref,nir_variable_mode mode)1555 nir_deref_mode_is(const nir_deref_instr *deref, nir_variable_mode mode)
1556 {
1557 assert(util_bitcount(mode) == 1 && (mode & nir_var_all));
1558 assert(deref->modes != 0);
1559
1560 /* This is only for "simple" cases so, if modes might interact with this
1561 * deref then the deref has to have a single mode.
1562 */
1563 if (nir_deref_mode_may_be(deref, mode)) {
1564 assert(util_bitcount(deref->modes) == 1);
1565 assert(deref->modes == mode);
1566 }
1567
1568 return deref->modes == mode;
1569 }
1570
1571 /** Returns true if deref has one of the given modes
1572 *
1573 * This returns true if the deref has exactly one possible mode and that mode
1574 * is one of the modes specified. If the deref may have one of those modes
1575 * but may also have a different mode (i.e. modes has multiple bits set), this
1576 * will assert-fail.
1577 */
1578 static inline bool
nir_deref_mode_is_one_of(const nir_deref_instr * deref,nir_variable_mode modes)1579 nir_deref_mode_is_one_of(const nir_deref_instr *deref, nir_variable_mode modes)
1580 {
1581 /* This is only for "simple" cases so, if modes might interact with this
1582 * deref then the deref has to have a single mode.
1583 */
1584 if (nir_deref_mode_may_be(deref, modes)) {
1585 assert(util_bitcount(deref->modes) == 1);
1586 assert(nir_deref_mode_must_be(deref, modes));
1587 }
1588
1589 return nir_deref_mode_may_be(deref, modes);
1590 }
1591
1592 /** Returns true if deref's possible modes lie in the given set of modes
1593 *
1594 * This returns true if the deref's modes lie in the given set of modes. If
1595 * the deref's modes overlap with the specified modes but aren't entirely
1596 * contained in the specified set of modes, this will assert-fail. In
1597 * particular, if this is used in a generic pointers scenario, the specified
1598 * modes has to contain all or none of the possible generic pointer modes.
1599 *
1600 * This is intended mostly for mass-lowering of derefs which might have
1601 * generic pointers.
1602 */
1603 static inline bool
nir_deref_mode_is_in_set(const nir_deref_instr * deref,nir_variable_mode modes)1604 nir_deref_mode_is_in_set(const nir_deref_instr *deref, nir_variable_mode modes)
1605 {
1606 if (nir_deref_mode_may_be(deref, modes))
1607 assert(nir_deref_mode_must_be(deref, modes));
1608
1609 return nir_deref_mode_may_be(deref, modes);
1610 }
1611
1612 static inline nir_deref_instr *nir_src_as_deref(nir_src src);
1613
1614 static inline nir_deref_instr *
nir_deref_instr_parent(const nir_deref_instr * instr)1615 nir_deref_instr_parent(const nir_deref_instr *instr)
1616 {
1617 if (instr->deref_type == nir_deref_type_var)
1618 return NULL;
1619 else
1620 return nir_src_as_deref(instr->parent);
1621 }
1622
1623 static inline nir_variable *
nir_deref_instr_get_variable(const nir_deref_instr * instr)1624 nir_deref_instr_get_variable(const nir_deref_instr *instr)
1625 {
1626 while (instr->deref_type != nir_deref_type_var) {
1627 if (instr->deref_type == nir_deref_type_cast)
1628 return NULL;
1629
1630 instr = nir_deref_instr_parent(instr);
1631 }
1632
1633 return instr->var;
1634 }
1635
1636 bool nir_deref_instr_has_indirect(nir_deref_instr *instr);
1637 bool nir_deref_instr_is_known_out_of_bounds(nir_deref_instr *instr);
1638 bool nir_deref_instr_has_complex_use(nir_deref_instr *instr);
1639
1640 bool nir_deref_instr_remove_if_unused(nir_deref_instr *instr);
1641
1642 unsigned nir_deref_instr_array_stride(nir_deref_instr *instr);
1643
1644 typedef struct {
1645 nir_instr instr;
1646
1647 struct nir_function *callee;
1648
1649 unsigned num_params;
1650 nir_src params[];
1651 } nir_call_instr;
1652
1653 #include "nir_intrinsics.h"
1654
1655 #define NIR_INTRINSIC_MAX_CONST_INDEX 7
1656
1657 /** Represents an intrinsic
1658 *
1659 * An intrinsic is an instruction type for handling things that are
1660 * more-or-less regular operations but don't just consume and produce SSA
1661 * values like ALU operations do. Intrinsics are not for things that have
1662 * special semantic meaning such as phi nodes and parallel copies.
1663 * Examples of intrinsics include variable load/store operations, system
1664 * value loads, and the like. Even though texturing more-or-less falls
1665 * under this category, texturing is its own instruction type because
1666 * trying to represent texturing with intrinsics would lead to a
1667 * combinatorial explosion of intrinsic opcodes.
1668 *
1669 * By having a single instruction type for handling a lot of different
1670 * cases, optimization passes can look for intrinsics and, for the most
1671 * part, completely ignore them. Each intrinsic type also has a few
1672 * possible flags that govern whether or not they can be reordered or
1673 * eliminated. That way passes like dead code elimination can still work
1674 * on intrisics without understanding the meaning of each.
1675 *
1676 * Each intrinsic has some number of constant indices, some number of
1677 * variables, and some number of sources. What these sources, variables,
1678 * and indices mean depends on the intrinsic and is documented with the
1679 * intrinsic declaration in nir_intrinsics.h. Intrinsics and texture
1680 * instructions are the only types of instruction that can operate on
1681 * variables.
1682 */
1683 typedef struct {
1684 nir_instr instr;
1685
1686 nir_intrinsic_op intrinsic;
1687
1688 nir_dest dest;
1689
1690 /** number of components if this is a vectorized intrinsic
1691 *
1692 * Similarly to ALU operations, some intrinsics are vectorized.
1693 * An intrinsic is vectorized if nir_intrinsic_infos.dest_components == 0.
1694 * For vectorized intrinsics, the num_components field specifies the
1695 * number of destination components and the number of source components
1696 * for all sources with nir_intrinsic_infos.src_components[i] == 0.
1697 */
1698 uint8_t num_components;
1699
1700 int const_index[NIR_INTRINSIC_MAX_CONST_INDEX];
1701
1702 nir_src src[];
1703 } nir_intrinsic_instr;
1704
1705 static inline nir_variable *
nir_intrinsic_get_var(nir_intrinsic_instr * intrin,unsigned i)1706 nir_intrinsic_get_var(nir_intrinsic_instr *intrin, unsigned i)
1707 {
1708 return nir_deref_instr_get_variable(nir_src_as_deref(intrin->src[i]));
1709 }
1710
1711 typedef enum {
1712 /* Memory ordering. */
1713 NIR_MEMORY_ACQUIRE = 1 << 0,
1714 NIR_MEMORY_RELEASE = 1 << 1,
1715 NIR_MEMORY_ACQ_REL = NIR_MEMORY_ACQUIRE | NIR_MEMORY_RELEASE,
1716
1717 /* Memory visibility operations. */
1718 NIR_MEMORY_MAKE_AVAILABLE = 1 << 2,
1719 NIR_MEMORY_MAKE_VISIBLE = 1 << 3,
1720 } nir_memory_semantics;
1721
1722 typedef enum {
1723 NIR_SCOPE_NONE,
1724 NIR_SCOPE_INVOCATION,
1725 NIR_SCOPE_SUBGROUP,
1726 NIR_SCOPE_SHADER_CALL,
1727 NIR_SCOPE_WORKGROUP,
1728 NIR_SCOPE_QUEUE_FAMILY,
1729 NIR_SCOPE_DEVICE,
1730 } nir_scope;
1731
1732 /**
1733 * \name NIR intrinsics semantic flags
1734 *
1735 * information about what the compiler can do with the intrinsics.
1736 *
1737 * \sa nir_intrinsic_info::flags
1738 */
1739 typedef enum {
1740 /**
1741 * whether the intrinsic can be safely eliminated if none of its output
1742 * value is not being used.
1743 */
1744 NIR_INTRINSIC_CAN_ELIMINATE = (1 << 0),
1745
1746 /**
1747 * Whether the intrinsic can be reordered with respect to any other
1748 * intrinsic, i.e. whether the only reordering dependencies of the
1749 * intrinsic are due to the register reads/writes.
1750 */
1751 NIR_INTRINSIC_CAN_REORDER = (1 << 1),
1752 } nir_intrinsic_semantic_flag;
1753
1754 /**
1755 * Maximum valid value for a nir align_mul value (in intrinsics or derefs).
1756 *
1757 * Offsets can be signed, so this is the largest power of two in int32_t.
1758 */
1759 #define NIR_ALIGN_MUL_MAX 0x40000000
1760
1761 typedef struct nir_io_semantics {
1762 unsigned location:7; /* gl_vert_attrib, gl_varying_slot, or gl_frag_result */
1763 unsigned num_slots:6; /* max 32, may be pessimistic with const indexing */
1764 unsigned dual_source_blend_index:1;
1765 unsigned fb_fetch_output:1; /* for GL_KHR_blend_equation_advanced */
1766 unsigned gs_streams:8; /* xxyyzzww: 2-bit stream index for each component */
1767 unsigned medium_precision:1; /* GLSL mediump qualifier */
1768 unsigned per_view:1;
1769 unsigned high_16bits:1; /* whether accessing low or high half of the slot */
1770 unsigned invariant:1; /* The variable has the invariant flag set */
1771 /* CLIP_DISTn, LAYER, VIEWPORT, and TESS_LEVEL_* have up to 3 uses:
1772 * - an output consumed by the next stage
1773 * - a system value output affecting fixed-func hardware, e.g. the clipper
1774 * - a transform feedback output written to memory
1775 * The following fields disable the first two. Transform feedback is disabled
1776 * by transform feedback info.
1777 */
1778 unsigned no_varying:1; /* whether this output isn't consumed by the next stage */
1779 unsigned no_sysval_output:1; /* whether this system value output has no
1780 effect due to current pipeline states */
1781 unsigned _pad:3;
1782 } nir_io_semantics;
1783
1784 /* Transform feedback info for 2 outputs. nir_intrinsic_store_output contains
1785 * this structure twice to support up to 4 outputs. The structure is limited
1786 * to 32 bits because it's stored in nir_intrinsic_instr::const_index[].
1787 */
1788 typedef struct nir_io_xfb {
1789 struct {
1790 /* start_component is equal to the index of out[]; add 2 for io_xfb2 */
1791 /* start_component is not relative to nir_intrinsic_component */
1792 /* get the stream index from nir_io_semantics */
1793 uint8_t num_components:4; /* max 4; if this is 0, xfb is disabled */
1794 uint8_t buffer:4; /* buffer index, max 3 */
1795 uint8_t offset; /* transform feedback buffer offset in dwords,
1796 max (1K - 4) bytes */
1797 } out[2];
1798 } nir_io_xfb;
1799
1800 unsigned
1801 nir_instr_xfb_write_mask(nir_intrinsic_instr *instr);
1802
1803 #define NIR_INTRINSIC_MAX_INPUTS 11
1804
1805 typedef struct {
1806 const char *name;
1807
1808 uint8_t num_srcs; /** < number of register/SSA inputs */
1809
1810 /** number of components of each input register
1811 *
1812 * If this value is 0, the number of components is given by the
1813 * num_components field of nir_intrinsic_instr. If this value is -1, the
1814 * intrinsic consumes however many components are provided and it is not
1815 * validated at all.
1816 */
1817 int8_t src_components[NIR_INTRINSIC_MAX_INPUTS];
1818
1819 bool has_dest;
1820
1821 /** number of components of the output register
1822 *
1823 * If this value is 0, the number of components is given by the
1824 * num_components field of nir_intrinsic_instr.
1825 */
1826 uint8_t dest_components;
1827
1828 /** bitfield of legal bit sizes */
1829 uint8_t dest_bit_sizes;
1830
1831 /** source which the destination bit size must match
1832 *
1833 * Some intrinsics, such as subgroup intrinsics, are data manipulation
1834 * intrinsics and they have similar bit-size rules to ALU ops. This enables
1835 * validation to validate a bit more and enables auto-generated builder code
1836 * to properly determine destination bit sizes automatically.
1837 */
1838 int8_t bit_size_src;
1839
1840 /** the number of constant indices used by the intrinsic */
1841 uint8_t num_indices;
1842
1843 /** list of indices */
1844 uint8_t indices[NIR_INTRINSIC_MAX_CONST_INDEX];
1845
1846 /** indicates the usage of intr->const_index[n] */
1847 uint8_t index_map[NIR_INTRINSIC_NUM_INDEX_FLAGS];
1848
1849 /** semantic flags for calls to this intrinsic */
1850 nir_intrinsic_semantic_flag flags;
1851 } nir_intrinsic_info;
1852
1853 extern const nir_intrinsic_info nir_intrinsic_infos[nir_num_intrinsics];
1854
1855 unsigned
1856 nir_intrinsic_src_components(const nir_intrinsic_instr *intr, unsigned srcn);
1857
1858 unsigned
1859 nir_intrinsic_dest_components(nir_intrinsic_instr *intr);
1860
1861 /**
1862 * Helper to copy const_index[] from src to dst, without assuming they
1863 * match in order.
1864 */
1865 void nir_intrinsic_copy_const_indices(nir_intrinsic_instr *dst, nir_intrinsic_instr *src);
1866
1867 #include "nir_intrinsics_indices.h"
1868
1869 static inline void
nir_intrinsic_set_align(nir_intrinsic_instr * intrin,unsigned align_mul,unsigned align_offset)1870 nir_intrinsic_set_align(nir_intrinsic_instr *intrin,
1871 unsigned align_mul, unsigned align_offset)
1872 {
1873 assert(util_is_power_of_two_nonzero(align_mul));
1874 assert(align_offset < align_mul);
1875 nir_intrinsic_set_align_mul(intrin, align_mul);
1876 nir_intrinsic_set_align_offset(intrin, align_offset);
1877 }
1878
1879 /** Returns a simple alignment for a load/store intrinsic offset
1880 *
1881 * Instead of the full mul+offset alignment scheme provided by the ALIGN_MUL
1882 * and ALIGN_OFFSET parameters, this helper takes both into account and
1883 * provides a single simple alignment parameter. The offset X is guaranteed
1884 * to satisfy X % align == 0.
1885 */
1886 static inline unsigned
nir_intrinsic_align(const nir_intrinsic_instr * intrin)1887 nir_intrinsic_align(const nir_intrinsic_instr *intrin)
1888 {
1889 const unsigned align_mul = nir_intrinsic_align_mul(intrin);
1890 const unsigned align_offset = nir_intrinsic_align_offset(intrin);
1891 assert(align_offset < align_mul);
1892 return align_offset ? 1 << (ffs(align_offset) - 1) : align_mul;
1893 }
1894
1895 static inline bool
nir_intrinsic_has_align(const nir_intrinsic_instr * intrin)1896 nir_intrinsic_has_align(const nir_intrinsic_instr *intrin)
1897 {
1898 return nir_intrinsic_has_align_mul(intrin) &&
1899 nir_intrinsic_has_align_offset(intrin);
1900 }
1901
1902 unsigned
1903 nir_image_intrinsic_coord_components(const nir_intrinsic_instr *instr);
1904
1905 /* Converts a image_deref_* intrinsic into a image_* one */
1906 void nir_rewrite_image_intrinsic(nir_intrinsic_instr *instr,
1907 nir_ssa_def *handle, bool bindless);
1908
1909 /* Determine if an intrinsic can be arbitrarily reordered and eliminated. */
1910 static inline bool
nir_intrinsic_can_reorder(nir_intrinsic_instr * instr)1911 nir_intrinsic_can_reorder(nir_intrinsic_instr *instr)
1912 {
1913 if (instr->intrinsic == nir_intrinsic_load_deref) {
1914 nir_deref_instr *deref = nir_src_as_deref(instr->src[0]);
1915 return nir_deref_mode_is_in_set(deref, nir_var_read_only_modes) ||
1916 (nir_intrinsic_access(instr) & ACCESS_CAN_REORDER);
1917 } else if (instr->intrinsic == nir_intrinsic_load_ssbo ||
1918 instr->intrinsic == nir_intrinsic_bindless_image_load ||
1919 instr->intrinsic == nir_intrinsic_image_deref_load ||
1920 instr->intrinsic == nir_intrinsic_image_load) {
1921 return nir_intrinsic_access(instr) & ACCESS_CAN_REORDER;
1922 } else {
1923 const nir_intrinsic_info *info =
1924 &nir_intrinsic_infos[instr->intrinsic];
1925 return (info->flags & NIR_INTRINSIC_CAN_ELIMINATE) &&
1926 (info->flags & NIR_INTRINSIC_CAN_REORDER);
1927 }
1928 }
1929
1930 bool nir_intrinsic_writes_external_memory(const nir_intrinsic_instr *instr);
1931
1932 /** Texture instruction source type */
1933 typedef enum {
1934 /** Texture coordinate
1935 *
1936 * Must have nir_tex_instr::coord_components components.
1937 */
1938 nir_tex_src_coord,
1939
1940 /** Projector
1941 *
1942 * The texture coordinate (except for the array component, if any) is
1943 * divided by this value before LOD computation and sampling.
1944 *
1945 * Must be a float scalar.
1946 */
1947 nir_tex_src_projector,
1948
1949 /** Shadow comparator
1950 *
1951 * For shadow sampling, the fetched texel values are compared against the
1952 * shadow comparator using the compare op specified by the sampler object
1953 * and converted to 1.0 if the comparison succeeds and 0.0 if it fails.
1954 * Interpolation happens after this conversion so the actual result may be
1955 * anywhere in the range [0.0, 1.0].
1956 *
1957 * Only valid if nir_tex_instr::is_shadow and must be a float scalar.
1958 */
1959 nir_tex_src_comparator,
1960
1961 /** Coordinate offset
1962 *
1963 * An integer value that is added to the texel address before sampling.
1964 * This is only allowed with operations that take an explicit LOD as it is
1965 * applied in integer texel space after LOD selection and not normalized
1966 * coordinate space.
1967 */
1968 nir_tex_src_offset,
1969
1970 /** LOD bias
1971 *
1972 * This value is added to the computed LOD before mip-mapping.
1973 */
1974 nir_tex_src_bias,
1975
1976 /** Explicit LOD */
1977 nir_tex_src_lod,
1978
1979 /** Min LOD
1980 *
1981 * The computed LOD is clamped to be at least as large as min_lod before
1982 * mip-mapping.
1983 */
1984 nir_tex_src_min_lod,
1985
1986 /** MSAA sample index */
1987 nir_tex_src_ms_index,
1988
1989 /** Intel-specific MSAA compression data */
1990 nir_tex_src_ms_mcs_intel,
1991
1992 /** Explicit horizontal (X-major) coordinate derivative */
1993 nir_tex_src_ddx,
1994
1995 /** Explicit vertical (Y-major) coordinate derivative */
1996 nir_tex_src_ddy,
1997
1998 /** Texture variable dereference */
1999 nir_tex_src_texture_deref,
2000
2001 /** Sampler variable dereference */
2002 nir_tex_src_sampler_deref,
2003
2004 /** Texture index offset
2005 *
2006 * This is added to nir_tex_instr::texture_index. Unless
2007 * nir_tex_instr::texture_non_uniform is set, this is guaranteed to be
2008 * dynamically uniform.
2009 */
2010 nir_tex_src_texture_offset,
2011
2012 /** Dynamically uniform sampler index offset
2013 *
2014 * This is added to nir_tex_instr::sampler_index. Unless
2015 * nir_tex_instr::sampler_non_uniform is set, this is guaranteed to be
2016 * dynamically uniform. This should not be present until GLSL ES 3.20, GLSL
2017 * 4.00, or ARB_gpu_shader5, because in ES 3.10 and GL 3.30 samplers said
2018 * "When aggregated into arrays within a shader, samplers can only be indexed
2019 * with a constant integral expression."
2020 */
2021 nir_tex_src_sampler_offset,
2022
2023 /** Bindless texture handle
2024 *
2025 * This is, unfortunately, a bit overloaded at the moment. There are
2026 * generally two types of bindless handles:
2027 *
2028 * 1. For GL_ARB_bindless bindless handles. These are part of the
2029 * GL/Gallium-level API and are always a 64-bit integer.
2030 *
2031 * 2. HW-specific handles. GL_ARB_bindless handles may be lowered to
2032 * these. Also, these are used by many Vulkan drivers to implement
2033 * descriptor sets, especially for UPDATE_AFTER_BIND descriptors.
2034 * The details of hardware handles (bit size, format, etc.) is
2035 * HW-specific.
2036 *
2037 * Because of this overloading and the resulting ambiguity, we currently
2038 * don't validate anything for these.
2039 */
2040 nir_tex_src_texture_handle,
2041
2042 /** Bindless sampler handle
2043 *
2044 * See nir_tex_src_texture_handle,
2045 */
2046 nir_tex_src_sampler_handle,
2047
2048 /** Plane index for multi-plane YCbCr textures */
2049 nir_tex_src_plane,
2050
2051 /**
2052 * Backend-specific vec4 tex src argument.
2053 *
2054 * Can be used to have NIR optimization (copy propagation, lower_vec_to_movs)
2055 * apply to the packing of the tex srcs. This lowering must only happen
2056 * after nir_lower_tex().
2057 *
2058 * The nir_tex_instr_src_type() of this argument is float, so no lowering
2059 * will happen if nir_lower_int_to_float is used.
2060 */
2061 nir_tex_src_backend1,
2062
2063 /** Second backend-specific vec4 tex src argument, see nir_tex_src_backend1. */
2064 nir_tex_src_backend2,
2065
2066 nir_num_tex_src_types
2067 } nir_tex_src_type;
2068
2069 /** A texture instruction source */
2070 typedef struct {
2071 /** Base source */
2072 nir_src src;
2073
2074 /** Type of this source */
2075 nir_tex_src_type src_type;
2076 } nir_tex_src;
2077
2078 /** Texture instruction opcode */
2079 typedef enum {
2080 nir_texop_tex, /**< Regular texture look-up */
2081 nir_texop_txb, /**< Texture look-up with LOD bias */
2082 nir_texop_txl, /**< Texture look-up with explicit LOD */
2083 nir_texop_txd, /**< Texture look-up with partial derivatives */
2084 nir_texop_txf, /**< Texel fetch with explicit LOD */
2085 nir_texop_txf_ms, /**< Multisample texture fetch */
2086 nir_texop_txf_ms_fb, /**< Multisample texture fetch from framebuffer */
2087 nir_texop_txf_ms_mcs_intel, /**< Multisample compression value fetch */
2088 nir_texop_txs, /**< Texture size */
2089 nir_texop_lod, /**< Texture lod query */
2090 nir_texop_tg4, /**< Texture gather */
2091 nir_texop_query_levels, /**< Texture levels query */
2092 nir_texop_texture_samples, /**< Texture samples query */
2093 nir_texop_samples_identical, /**< Query whether all samples are definitely
2094 * identical.
2095 */
2096 nir_texop_tex_prefetch, /**< Regular texture look-up, eligible for pre-dispatch */
2097 nir_texop_fragment_fetch_amd, /**< Multisample fragment color texture fetch */
2098 nir_texop_fragment_mask_fetch_amd, /**< Multisample fragment mask texture fetch */
2099 } nir_texop;
2100
2101 /** Represents a texture instruction */
2102 typedef struct {
2103 /** Base instruction */
2104 nir_instr instr;
2105
2106 /** Dimensionality of the texture operation
2107 *
2108 * This will typically match the dimensionality of the texture deref type
2109 * if a nir_tex_src_texture_deref is present. However, it may not if
2110 * texture lowering has occurred.
2111 */
2112 enum glsl_sampler_dim sampler_dim;
2113
2114 /** ALU type of the destination
2115 *
2116 * This is the canonical sampled type for this texture operation and may
2117 * not exactly match the sampled type of the deref type when a
2118 * nir_tex_src_texture_deref is present. For OpenCL, the sampled type of
2119 * the texture deref will be GLSL_TYPE_VOID and this is allowed to be
2120 * anything. With SPIR-V, the signedness of integer types is allowed to
2121 * differ. For all APIs, the bit size may differ if the driver has done
2122 * any sort of mediump or similar lowering since texture types always have
2123 * 32-bit sampled types.
2124 */
2125 nir_alu_type dest_type;
2126
2127 /** Texture opcode */
2128 nir_texop op;
2129
2130 /** Destination */
2131 nir_dest dest;
2132
2133 /** Array of sources
2134 *
2135 * This array has nir_tex_instr::num_srcs elements
2136 */
2137 nir_tex_src *src;
2138
2139 /** Number of sources */
2140 unsigned num_srcs;
2141
2142 /** Number of components in the coordinate, if any */
2143 unsigned coord_components;
2144
2145 /** True if the texture instruction acts on an array texture */
2146 bool is_array;
2147
2148 /** True if the texture instruction performs a shadow comparison
2149 *
2150 * If this is true, the texture instruction must have a
2151 * nir_tex_src_comparator.
2152 */
2153 bool is_shadow;
2154
2155 /**
2156 * If is_shadow is true, whether this is the old-style shadow that outputs
2157 * 4 components or the new-style shadow that outputs 1 component.
2158 */
2159 bool is_new_style_shadow;
2160
2161 /**
2162 * True if this texture instruction should return a sparse residency code.
2163 * The code is in the last component of the result.
2164 */
2165 bool is_sparse;
2166
2167 /** nir_texop_tg4 component selector
2168 *
2169 * This determines which RGBA component is gathered.
2170 */
2171 unsigned component : 2;
2172
2173 /** Validation needs to know this for gradient component count */
2174 unsigned array_is_lowered_cube : 1;
2175
2176 /** Gather offsets */
2177 int8_t tg4_offsets[4][2];
2178
2179 /** True if the texture index or handle is not dynamically uniform */
2180 bool texture_non_uniform;
2181
2182 /** True if the sampler index or handle is not dynamically uniform.
2183 *
2184 * This may be set when VK_EXT_descriptor_indexing is supported and the
2185 * appropriate capability is enabled.
2186 *
2187 * This should always be false in GLSL (GLSL ES 3.20 says "When aggregated
2188 * into arrays within a shader, opaque types can only be indexed with a
2189 * dynamically uniform integral expression", and GLSL 4.60 says "When
2190 * aggregated into arrays within a shader, [texture, sampler, and
2191 * samplerShadow] types can only be indexed with a dynamically uniform
2192 * expression, or texture lookup will result in undefined values.").
2193 */
2194 bool sampler_non_uniform;
2195
2196 /** The texture index
2197 *
2198 * If this texture instruction has a nir_tex_src_texture_offset source,
2199 * then the texture index is given by texture_index + texture_offset.
2200 */
2201 unsigned texture_index;
2202
2203 /** The sampler index
2204 *
2205 * The following operations do not require a sampler and, as such, this
2206 * field should be ignored:
2207 * - nir_texop_txf
2208 * - nir_texop_txf_ms
2209 * - nir_texop_txs
2210 * - nir_texop_query_levels
2211 * - nir_texop_texture_samples
2212 * - nir_texop_samples_identical
2213 *
2214 * If this texture instruction has a nir_tex_src_sampler_offset source,
2215 * then the sampler index is given by sampler_index + sampler_offset.
2216 */
2217 unsigned sampler_index;
2218 } nir_tex_instr;
2219
2220 /**
2221 * Returns true if the texture operation requires a sampler as a general rule
2222 *
2223 * Note that the specific hw/driver backend could require to a sampler
2224 * object/configuration packet in any case, for some other reason.
2225 *
2226 * @see nir_tex_instr::sampler_index.
2227 */
2228 bool nir_tex_instr_need_sampler(const nir_tex_instr *instr);
2229
2230 /** Returns the number of components returned by this nir_tex_instr
2231 *
2232 * Useful for code building texture instructions when you don't want to think
2233 * about how many components a particular texture op returns. This does not
2234 * include the sparse residency code.
2235 */
2236 unsigned
2237 nir_tex_instr_result_size(const nir_tex_instr *instr);
2238
2239 /**
2240 * Returns the destination size of this nir_tex_instr including the sparse
2241 * residency code, if any.
2242 */
2243 static inline unsigned
nir_tex_instr_dest_size(const nir_tex_instr * instr)2244 nir_tex_instr_dest_size(const nir_tex_instr *instr)
2245 {
2246 /* One more component is needed for the residency code. */
2247 return nir_tex_instr_result_size(instr) + instr->is_sparse;
2248 }
2249
2250 /**
2251 * Returns true if this texture operation queries something about the texture
2252 * rather than actually sampling it.
2253 */
2254 bool
2255 nir_tex_instr_is_query(const nir_tex_instr *instr);
2256
2257 /** Returns true if this texture instruction does implicit derivatives
2258 *
2259 * This is important as there are extra control-flow rules around derivatives
2260 * and texture instructions which perform them implicitly.
2261 */
2262 bool
2263 nir_tex_instr_has_implicit_derivative(const nir_tex_instr *instr);
2264
2265 /** Returns the ALU type of the given texture instruction source */
2266 nir_alu_type
2267 nir_tex_instr_src_type(const nir_tex_instr *instr, unsigned src);
2268
2269 /**
2270 * Returns the number of components required by the given texture instruction
2271 * source
2272 */
2273 unsigned
2274 nir_tex_instr_src_size(const nir_tex_instr *instr, unsigned src);
2275
2276 /**
2277 * Returns the index of the texture instruction source with the given
2278 * nir_tex_src_type or -1 if no such source exists.
2279 */
2280 static inline int
nir_tex_instr_src_index(const nir_tex_instr * instr,nir_tex_src_type type)2281 nir_tex_instr_src_index(const nir_tex_instr *instr, nir_tex_src_type type)
2282 {
2283 for (unsigned i = 0; i < instr->num_srcs; i++)
2284 if (instr->src[i].src_type == type)
2285 return (int) i;
2286
2287 return -1;
2288 }
2289
2290 /** Adds a source to a texture instruction */
2291 void nir_tex_instr_add_src(nir_tex_instr *tex,
2292 nir_tex_src_type src_type,
2293 nir_src src);
2294
2295 /** Removes a source from a texture instruction */
2296 void nir_tex_instr_remove_src(nir_tex_instr *tex, unsigned src_idx);
2297
2298 bool nir_tex_instr_has_explicit_tg4_offsets(nir_tex_instr *tex);
2299
2300 typedef struct {
2301 nir_instr instr;
2302
2303 nir_ssa_def def;
2304
2305 nir_const_value value[];
2306 } nir_load_const_instr;
2307
2308 typedef enum {
2309 /** Return from a function
2310 *
2311 * This instruction is a classic function return. It jumps to
2312 * nir_function_impl::end_block. No return value is provided in this
2313 * instruction. Instead, the function is expected to write any return
2314 * data to a deref passed in from the caller.
2315 */
2316 nir_jump_return,
2317
2318 /** Immediately exit the current shader
2319 *
2320 * This instruction is roughly the equivalent of C's "exit()" in that it
2321 * immediately terminates the current shader invocation. From a CFG
2322 * perspective, it looks like a jump to nir_function_impl::end_block but
2323 * it actually jumps to the end block of the shader entrypoint. A halt
2324 * instruction in the shader entrypoint itself is semantically identical
2325 * to a return.
2326 *
2327 * For shaders with built-in I/O, any outputs written prior to a halt
2328 * instruction remain written and any outputs not written prior to the
2329 * halt have undefined values. It does NOT cause an implicit discard of
2330 * written results. If one wants discard results in a fragment shader,
2331 * for instance, a discard or demote intrinsic is required.
2332 */
2333 nir_jump_halt,
2334
2335 /** Break out of the inner-most loop
2336 *
2337 * This has the same semantics as C's "break" statement.
2338 */
2339 nir_jump_break,
2340
2341 /** Jump back to the top of the inner-most loop
2342 *
2343 * This has the same semantics as C's "continue" statement assuming that a
2344 * NIR loop is implemented as "while (1) { body }".
2345 */
2346 nir_jump_continue,
2347
2348 /** Jumps for unstructured CFG.
2349 *
2350 * As within an unstructured CFG we can't rely on block ordering we need to
2351 * place explicit jumps at the end of every block.
2352 */
2353 nir_jump_goto,
2354 nir_jump_goto_if,
2355 } nir_jump_type;
2356
2357 typedef struct {
2358 nir_instr instr;
2359 nir_jump_type type;
2360 nir_src condition;
2361 struct nir_block *target;
2362 struct nir_block *else_target;
2363 } nir_jump_instr;
2364
2365 /* creates a new SSA variable in an undefined state */
2366
2367 typedef struct {
2368 nir_instr instr;
2369 nir_ssa_def def;
2370 } nir_ssa_undef_instr;
2371
2372 typedef struct {
2373 struct exec_node node;
2374
2375 /* The predecessor block corresponding to this source */
2376 struct nir_block *pred;
2377
2378 nir_src src;
2379 } nir_phi_src;
2380
2381 #define nir_foreach_phi_src(phi_src, phi) \
2382 foreach_list_typed(nir_phi_src, phi_src, node, &(phi)->srcs)
2383 #define nir_foreach_phi_src_safe(phi_src, phi) \
2384 foreach_list_typed_safe(nir_phi_src, phi_src, node, &(phi)->srcs)
2385
2386 typedef struct {
2387 nir_instr instr;
2388
2389 struct exec_list srcs; /** < list of nir_phi_src */
2390
2391 nir_dest dest;
2392 } nir_phi_instr;
2393
2394 static inline nir_phi_src *
nir_phi_get_src_from_block(nir_phi_instr * phi,struct nir_block * block)2395 nir_phi_get_src_from_block(nir_phi_instr *phi, struct nir_block *block)
2396 {
2397 nir_foreach_phi_src(src, phi) {
2398 if (src->pred == block)
2399 return src;
2400 }
2401
2402 assert(!"Block is not a predecessor of phi.");
2403 return NULL;
2404 }
2405
2406 typedef struct {
2407 struct exec_node node;
2408 nir_src src;
2409 nir_dest dest;
2410 } nir_parallel_copy_entry;
2411
2412 #define nir_foreach_parallel_copy_entry(entry, pcopy) \
2413 foreach_list_typed(nir_parallel_copy_entry, entry, node, &(pcopy)->entries)
2414
2415 typedef struct {
2416 nir_instr instr;
2417
2418 /* A list of nir_parallel_copy_entrys. The sources of all of the
2419 * entries are copied to the corresponding destinations "in parallel".
2420 * In other words, if we have two entries: a -> b and b -> a, the values
2421 * get swapped.
2422 */
2423 struct exec_list entries;
2424 } nir_parallel_copy_instr;
2425
2426 NIR_DEFINE_CAST(nir_instr_as_alu, nir_instr, nir_alu_instr, instr,
2427 type, nir_instr_type_alu)
2428 NIR_DEFINE_CAST(nir_instr_as_deref, nir_instr, nir_deref_instr, instr,
2429 type, nir_instr_type_deref)
2430 NIR_DEFINE_CAST(nir_instr_as_call, nir_instr, nir_call_instr, instr,
2431 type, nir_instr_type_call)
2432 NIR_DEFINE_CAST(nir_instr_as_jump, nir_instr, nir_jump_instr, instr,
2433 type, nir_instr_type_jump)
2434 NIR_DEFINE_CAST(nir_instr_as_tex, nir_instr, nir_tex_instr, instr,
2435 type, nir_instr_type_tex)
2436 NIR_DEFINE_CAST(nir_instr_as_intrinsic, nir_instr, nir_intrinsic_instr, instr,
2437 type, nir_instr_type_intrinsic)
2438 NIR_DEFINE_CAST(nir_instr_as_load_const, nir_instr, nir_load_const_instr, instr,
2439 type, nir_instr_type_load_const)
2440 NIR_DEFINE_CAST(nir_instr_as_ssa_undef, nir_instr, nir_ssa_undef_instr, instr,
2441 type, nir_instr_type_ssa_undef)
2442 NIR_DEFINE_CAST(nir_instr_as_phi, nir_instr, nir_phi_instr, instr,
2443 type, nir_instr_type_phi)
2444 NIR_DEFINE_CAST(nir_instr_as_parallel_copy, nir_instr,
2445 nir_parallel_copy_instr, instr,
2446 type, nir_instr_type_parallel_copy)
2447
2448
2449 #define NIR_DEFINE_SRC_AS_CONST(type, suffix) \
2450 static inline type \
2451 nir_src_comp_as_##suffix(nir_src src, unsigned comp) \
2452 { \
2453 assert(nir_src_is_const(src)); \
2454 nir_load_const_instr *load = \
2455 nir_instr_as_load_const(src.ssa->parent_instr); \
2456 assert(comp < load->def.num_components); \
2457 return nir_const_value_as_##suffix(load->value[comp], \
2458 load->def.bit_size); \
2459 } \
2460 \
2461 static inline type \
2462 nir_src_as_##suffix(nir_src src) \
2463 { \
2464 assert(nir_src_num_components(src) == 1); \
2465 return nir_src_comp_as_##suffix(src, 0); \
2466 }
2467
2468 NIR_DEFINE_SRC_AS_CONST(int64_t, int)
2469 NIR_DEFINE_SRC_AS_CONST(uint64_t, uint)
2470 NIR_DEFINE_SRC_AS_CONST(bool, bool)
2471 NIR_DEFINE_SRC_AS_CONST(double, float)
2472
2473 #undef NIR_DEFINE_SRC_AS_CONST
2474
2475
2476 typedef struct {
2477 nir_ssa_def *def;
2478 unsigned comp;
2479 } nir_ssa_scalar;
2480
2481 static inline bool
nir_ssa_scalar_is_const(nir_ssa_scalar s)2482 nir_ssa_scalar_is_const(nir_ssa_scalar s)
2483 {
2484 return s.def->parent_instr->type == nir_instr_type_load_const;
2485 }
2486
2487 static inline nir_const_value
nir_ssa_scalar_as_const_value(nir_ssa_scalar s)2488 nir_ssa_scalar_as_const_value(nir_ssa_scalar s)
2489 {
2490 assert(s.comp < s.def->num_components);
2491 nir_load_const_instr *load = nir_instr_as_load_const(s.def->parent_instr);
2492 return load->value[s.comp];
2493 }
2494
2495 #define NIR_DEFINE_SCALAR_AS_CONST(type, suffix) \
2496 static inline type \
2497 nir_ssa_scalar_as_##suffix(nir_ssa_scalar s) \
2498 { \
2499 return nir_const_value_as_##suffix( \
2500 nir_ssa_scalar_as_const_value(s), s.def->bit_size); \
2501 }
2502
NIR_DEFINE_SCALAR_AS_CONST(int64_t,int)2503 NIR_DEFINE_SCALAR_AS_CONST(int64_t, int)
2504 NIR_DEFINE_SCALAR_AS_CONST(uint64_t, uint)
2505 NIR_DEFINE_SCALAR_AS_CONST(bool, bool)
2506 NIR_DEFINE_SCALAR_AS_CONST(double, float)
2507
2508 #undef NIR_DEFINE_SCALAR_AS_CONST
2509
2510 static inline bool
2511 nir_ssa_scalar_is_alu(nir_ssa_scalar s)
2512 {
2513 return s.def->parent_instr->type == nir_instr_type_alu;
2514 }
2515
2516 static inline nir_op
nir_ssa_scalar_alu_op(nir_ssa_scalar s)2517 nir_ssa_scalar_alu_op(nir_ssa_scalar s)
2518 {
2519 return nir_instr_as_alu(s.def->parent_instr)->op;
2520 }
2521
2522 static inline nir_ssa_scalar
nir_ssa_scalar_chase_alu_src(nir_ssa_scalar s,unsigned alu_src_idx)2523 nir_ssa_scalar_chase_alu_src(nir_ssa_scalar s, unsigned alu_src_idx)
2524 {
2525 nir_ssa_scalar out = { NULL, 0 };
2526
2527 nir_alu_instr *alu = nir_instr_as_alu(s.def->parent_instr);
2528 assert(alu_src_idx < nir_op_infos[alu->op].num_inputs);
2529
2530 /* Our component must be written */
2531 assert(s.comp < s.def->num_components);
2532 assert(alu->dest.write_mask & (1u << s.comp));
2533
2534 assert(alu->src[alu_src_idx].src.is_ssa);
2535 out.def = alu->src[alu_src_idx].src.ssa;
2536
2537 if (nir_op_infos[alu->op].input_sizes[alu_src_idx] == 0) {
2538 /* The ALU src is unsized so the source component follows the
2539 * destination component.
2540 */
2541 out.comp = alu->src[alu_src_idx].swizzle[s.comp];
2542 } else {
2543 /* This is a sized source so all source components work together to
2544 * produce all the destination components. Since we need to return a
2545 * scalar, this only works if the source is a scalar.
2546 */
2547 assert(nir_op_infos[alu->op].input_sizes[alu_src_idx] == 1);
2548 out.comp = alu->src[alu_src_idx].swizzle[0];
2549 }
2550 assert(out.comp < out.def->num_components);
2551
2552 return out;
2553 }
2554
2555 nir_ssa_scalar nir_ssa_scalar_chase_movs(nir_ssa_scalar s);
2556
2557 static inline nir_ssa_scalar
nir_get_ssa_scalar(nir_ssa_def * def,unsigned channel)2558 nir_get_ssa_scalar(nir_ssa_def *def, unsigned channel)
2559 {
2560 nir_ssa_scalar s = { def, channel };
2561 return s;
2562 }
2563
2564 /** Returns a nir_ssa_scalar where we've followed the bit-exact mov/vec use chain to the original definition */
2565 static inline nir_ssa_scalar
nir_ssa_scalar_resolved(nir_ssa_def * def,unsigned channel)2566 nir_ssa_scalar_resolved(nir_ssa_def *def, unsigned channel)
2567 {
2568 return nir_ssa_scalar_chase_movs(nir_get_ssa_scalar(def, channel));
2569 }
2570
2571
2572 typedef struct {
2573 bool success;
2574
2575 nir_variable *var;
2576 unsigned desc_set;
2577 unsigned binding;
2578 unsigned num_indices;
2579 nir_src indices[4];
2580 bool read_first_invocation;
2581 } nir_binding;
2582
2583 nir_binding nir_chase_binding(nir_src rsrc);
2584 nir_variable *nir_get_binding_variable(struct nir_shader *shader, nir_binding binding);
2585
2586
2587 /*
2588 * Control flow
2589 *
2590 * Control flow consists of a tree of control flow nodes, which include
2591 * if-statements and loops. The leaves of the tree are basic blocks, lists of
2592 * instructions that always run start-to-finish. Each basic block also keeps
2593 * track of its successors (blocks which may run immediately after the current
2594 * block) and predecessors (blocks which could have run immediately before the
2595 * current block). Each function also has a start block and an end block which
2596 * all return statements point to (which is always empty). Together, all the
2597 * blocks with their predecessors and successors make up the control flow
2598 * graph (CFG) of the function. There are helpers that modify the tree of
2599 * control flow nodes while modifying the CFG appropriately; these should be
2600 * used instead of modifying the tree directly.
2601 */
2602
2603 typedef enum {
2604 nir_cf_node_block,
2605 nir_cf_node_if,
2606 nir_cf_node_loop,
2607 nir_cf_node_function
2608 } nir_cf_node_type;
2609
2610 typedef struct nir_cf_node {
2611 struct exec_node node;
2612 nir_cf_node_type type;
2613 struct nir_cf_node *parent;
2614 } nir_cf_node;
2615
2616 typedef struct nir_block {
2617 nir_cf_node cf_node;
2618
2619 struct exec_list instr_list; /** < list of nir_instr */
2620
2621 /** generic block index; generated by nir_index_blocks */
2622 unsigned index;
2623
2624 /*
2625 * Each block can only have up to 2 successors, so we put them in a simple
2626 * array - no need for anything more complicated.
2627 */
2628 struct nir_block *successors[2];
2629
2630 /* Set of nir_block predecessors in the CFG */
2631 struct set *predecessors;
2632
2633 /*
2634 * this node's immediate dominator in the dominance tree - set to NULL for
2635 * the start block.
2636 */
2637 struct nir_block *imm_dom;
2638
2639 /* This node's children in the dominance tree */
2640 unsigned num_dom_children;
2641 struct nir_block **dom_children;
2642
2643 /* Set of nir_blocks on the dominance frontier of this block */
2644 struct set *dom_frontier;
2645
2646 /*
2647 * These two indices have the property that dom_{pre,post}_index for each
2648 * child of this block in the dominance tree will always be between
2649 * dom_pre_index and dom_post_index for this block, which makes testing if
2650 * a given block is dominated by another block an O(1) operation.
2651 */
2652 uint32_t dom_pre_index, dom_post_index;
2653
2654 /**
2655 * Value just before the first nir_instr->index in the block, but after
2656 * end_ip that of any predecessor block.
2657 */
2658 uint32_t start_ip;
2659 /**
2660 * Value just after the last nir_instr->index in the block, but before the
2661 * start_ip of any successor block.
2662 */
2663 uint32_t end_ip;
2664
2665 /* SSA def live in and out for this block; used for liveness analysis.
2666 * Indexed by ssa_def->index
2667 */
2668 BITSET_WORD *live_in;
2669 BITSET_WORD *live_out;
2670 } nir_block;
2671
2672 static inline bool
nir_block_is_reachable(nir_block * b)2673 nir_block_is_reachable(nir_block *b)
2674 {
2675 /* See also nir_block_dominates */
2676 return b->dom_post_index != 0;
2677 }
2678
2679 static inline nir_instr *
nir_block_first_instr(nir_block * block)2680 nir_block_first_instr(nir_block *block)
2681 {
2682 struct exec_node *head = exec_list_get_head(&block->instr_list);
2683 return exec_node_data(nir_instr, head, node);
2684 }
2685
2686 static inline nir_instr *
nir_block_last_instr(nir_block * block)2687 nir_block_last_instr(nir_block *block)
2688 {
2689 struct exec_node *tail = exec_list_get_tail(&block->instr_list);
2690 return exec_node_data(nir_instr, tail, node);
2691 }
2692
2693 static inline bool
nir_block_ends_in_jump(nir_block * block)2694 nir_block_ends_in_jump(nir_block *block)
2695 {
2696 return !exec_list_is_empty(&block->instr_list) &&
2697 nir_block_last_instr(block)->type == nir_instr_type_jump;
2698 }
2699
2700 static inline bool
nir_block_ends_in_return_or_halt(nir_block * block)2701 nir_block_ends_in_return_or_halt(nir_block *block)
2702 {
2703 if (exec_list_is_empty(&block->instr_list))
2704 return false;
2705
2706 nir_instr *instr = nir_block_last_instr(block);
2707 if (instr->type != nir_instr_type_jump)
2708 return false;
2709
2710 nir_jump_instr *jump_instr = nir_instr_as_jump(instr);
2711 return jump_instr->type == nir_jump_return ||
2712 jump_instr->type == nir_jump_halt;
2713 }
2714
2715 static inline bool
nir_block_ends_in_break(nir_block * block)2716 nir_block_ends_in_break(nir_block *block)
2717 {
2718 if (exec_list_is_empty(&block->instr_list))
2719 return false;
2720
2721 nir_instr *instr = nir_block_last_instr(block);
2722 return instr->type == nir_instr_type_jump &&
2723 nir_instr_as_jump(instr)->type == nir_jump_break;
2724 }
2725
2726 #define nir_foreach_instr(instr, block) \
2727 foreach_list_typed(nir_instr, instr, node, &(block)->instr_list)
2728 #define nir_foreach_instr_reverse(instr, block) \
2729 foreach_list_typed_reverse(nir_instr, instr, node, &(block)->instr_list)
2730 #define nir_foreach_instr_safe(instr, block) \
2731 foreach_list_typed_safe(nir_instr, instr, node, &(block)->instr_list)
2732 #define nir_foreach_instr_reverse_safe(instr, block) \
2733 foreach_list_typed_reverse_safe(nir_instr, instr, node, &(block)->instr_list)
2734
2735 static inline nir_phi_instr *
nir_block_last_phi_instr(nir_block * block)2736 nir_block_last_phi_instr(nir_block *block)
2737 {
2738 nir_phi_instr *last_phi = NULL;
2739 nir_foreach_instr(instr, block) {
2740 if (instr->type == nir_instr_type_phi)
2741 last_phi = nir_instr_as_phi(instr);
2742 else
2743 return last_phi;
2744 }
2745 return last_phi;
2746 }
2747
2748 typedef enum {
2749 nir_selection_control_none = 0x0,
2750 nir_selection_control_flatten = 0x1,
2751 nir_selection_control_dont_flatten = 0x2,
2752 } nir_selection_control;
2753
2754 typedef struct nir_if {
2755 nir_cf_node cf_node;
2756 nir_src condition;
2757 nir_selection_control control;
2758
2759 struct exec_list then_list; /** < list of nir_cf_node */
2760 struct exec_list else_list; /** < list of nir_cf_node */
2761 } nir_if;
2762
2763 typedef struct {
2764 nir_if *nif;
2765
2766 /** Instruction that generates nif::condition. */
2767 nir_instr *conditional_instr;
2768
2769 /** Block within ::nif that has the break instruction. */
2770 nir_block *break_block;
2771
2772 /** Last block for the then- or else-path that does not contain the break. */
2773 nir_block *continue_from_block;
2774
2775 /** True when ::break_block is in the else-path of ::nif. */
2776 bool continue_from_then;
2777 bool induction_rhs;
2778
2779 /* This is true if the terminators exact trip count is unknown. For
2780 * example:
2781 *
2782 * for (int i = 0; i < imin(x, 4); i++)
2783 * ...
2784 *
2785 * Here loop analysis would have set a max_trip_count of 4 however we dont
2786 * know for sure that this is the exact trip count.
2787 */
2788 bool exact_trip_count_unknown;
2789
2790 struct list_head loop_terminator_link;
2791 } nir_loop_terminator;
2792
2793 typedef struct {
2794 /* Induction variable. */
2795 nir_ssa_def *def;
2796
2797 /* Init statement with only uniform. */
2798 nir_src *init_src;
2799
2800 /* Update statement with only uniform. */
2801 nir_alu_src *update_src;
2802 } nir_loop_induction_variable;
2803
2804 typedef struct {
2805 /* Estimated cost (in number of instructions) of the loop */
2806 unsigned instr_cost;
2807
2808 /* Guessed trip count based on array indexing */
2809 unsigned guessed_trip_count;
2810
2811 /* Maximum number of times the loop is run (if known) */
2812 unsigned max_trip_count;
2813
2814 /* Do we know the exact number of times the loop will be run */
2815 bool exact_trip_count_known;
2816
2817 /* Unroll the loop regardless of its size */
2818 bool force_unroll;
2819
2820 /* Does the loop contain complex loop terminators, continues or other
2821 * complex behaviours? If this is true we can't rely on
2822 * loop_terminator_list to be complete or accurate.
2823 */
2824 bool complex_loop;
2825
2826 nir_loop_terminator *limiting_terminator;
2827
2828 /* A list of loop_terminators terminating this loop. */
2829 struct list_head loop_terminator_list;
2830
2831 /* array of induction variables for this loop */
2832 nir_loop_induction_variable *induction_vars;
2833 unsigned num_induction_vars;
2834 } nir_loop_info;
2835
2836 typedef enum {
2837 nir_loop_control_none = 0x0,
2838 nir_loop_control_unroll = 0x1,
2839 nir_loop_control_dont_unroll = 0x2,
2840 } nir_loop_control;
2841
2842 typedef struct {
2843 nir_cf_node cf_node;
2844
2845 struct exec_list body; /** < list of nir_cf_node */
2846
2847 nir_loop_info *info;
2848 nir_loop_control control;
2849 bool partially_unrolled;
2850 bool divergent;
2851 } nir_loop;
2852
2853 /**
2854 * Various bits of metadata that can may be created or required by
2855 * optimization and analysis passes
2856 */
2857 typedef enum {
2858 nir_metadata_none = 0x0,
2859
2860 /** Indicates that nir_block::index values are valid.
2861 *
2862 * The start block has index 0 and they increase through a natural walk of
2863 * the CFG. nir_function_impl::num_blocks is the number of blocks and
2864 * every block index is in the range [0, nir_function_impl::num_blocks].
2865 *
2866 * A pass can preserve this metadata type if it doesn't touch the CFG.
2867 */
2868 nir_metadata_block_index = 0x1,
2869
2870 /** Indicates that block dominance information is valid
2871 *
2872 * This includes:
2873 *
2874 * - nir_block::num_dom_children
2875 * - nir_block::dom_children
2876 * - nir_block::dom_frontier
2877 * - nir_block::dom_pre_index
2878 * - nir_block::dom_post_index
2879 *
2880 * A pass can preserve this metadata type if it doesn't touch the CFG.
2881 */
2882 nir_metadata_dominance = 0x2,
2883
2884 /** Indicates that SSA def data-flow liveness information is valid
2885 *
2886 * This includes:
2887 *
2888 * - nir_block::live_in
2889 * - nir_block::live_out
2890 *
2891 * A pass can preserve this metadata type if it never adds or removes any
2892 * SSA defs or uses of SSA defs (most passes shouldn't preserve this
2893 * metadata type).
2894 */
2895 nir_metadata_live_ssa_defs = 0x4,
2896
2897 /** A dummy metadata value to track when a pass forgot to call
2898 * nir_metadata_preserve.
2899 *
2900 * A pass should always clear this value even if it doesn't make any
2901 * progress to indicate that it thought about preserving metadata.
2902 */
2903 nir_metadata_not_properly_reset = 0x8,
2904
2905 /** Indicates that loop analysis information is valid.
2906 *
2907 * This includes everything pointed to by nir_loop::info.
2908 *
2909 * A pass can preserve this metadata type if it is guaranteed to not affect
2910 * any loop metadata. However, since loop metadata includes things like
2911 * loop counts which depend on arithmetic in the loop, this is very hard to
2912 * determine. Most passes shouldn't preserve this metadata type.
2913 */
2914 nir_metadata_loop_analysis = 0x10,
2915
2916 /** Indicates that nir_instr::index values are valid.
2917 *
2918 * The start instruction has index 0 and they increase through a natural
2919 * walk of instructions in blocks in the CFG. The indices my have holes
2920 * after passes such as DCE.
2921 *
2922 * A pass can preserve this metadata type if it never adds or moves any
2923 * instructions (most passes shouldn't preserve this metadata type), but
2924 * can preserve it if it only removes instructions.
2925 */
2926 nir_metadata_instr_index = 0x20,
2927
2928 /** All metadata
2929 *
2930 * This includes all nir_metadata flags except not_properly_reset. Passes
2931 * which do not change the shader in any way should call
2932 *
2933 * nir_metadata_preserve(impl, nir_metadata_all);
2934 */
2935 nir_metadata_all = ~nir_metadata_not_properly_reset,
2936 } nir_metadata;
2937 MESA_DEFINE_CPP_ENUM_BITFIELD_OPERATORS(nir_metadata)
2938
2939 typedef struct {
2940 nir_cf_node cf_node;
2941
2942 /** pointer to the function of which this is an implementation */
2943 struct nir_function *function;
2944
2945 /**
2946 * For entrypoints, a pointer to a nir_function_impl which runs before
2947 * it, once per draw or dispatch, communicating via store_preamble and
2948 * load_preamble intrinsics. If NULL then there is no preamble.
2949 */
2950 struct nir_function *preamble;
2951
2952 struct exec_list body; /** < list of nir_cf_node */
2953
2954 nir_block *end_block;
2955
2956 /** list for all local variables in the function */
2957 struct exec_list locals;
2958
2959 /** list of local registers in the function */
2960 struct exec_list registers;
2961
2962 /** next available local register index */
2963 unsigned reg_alloc;
2964
2965 /** next available SSA value index */
2966 unsigned ssa_alloc;
2967
2968 /* total number of basic blocks, only valid when block_index_dirty = false */
2969 unsigned num_blocks;
2970
2971 /** True if this nir_function_impl uses structured control-flow
2972 *
2973 * Structured nir_function_impls have different validation rules.
2974 */
2975 bool structured;
2976
2977 nir_metadata valid_metadata;
2978 } nir_function_impl;
2979
2980 #define nir_foreach_function_temp_variable(var, impl) \
2981 foreach_list_typed(nir_variable, var, node, &(impl)->locals)
2982
2983 #define nir_foreach_function_temp_variable_safe(var, impl) \
2984 foreach_list_typed_safe(nir_variable, var, node, &(impl)->locals)
2985
2986 ATTRIBUTE_RETURNS_NONNULL static inline nir_block *
nir_start_block(nir_function_impl * impl)2987 nir_start_block(nir_function_impl *impl)
2988 {
2989 return (nir_block *) impl->body.head_sentinel.next;
2990 }
2991
2992 ATTRIBUTE_RETURNS_NONNULL static inline nir_block *
nir_impl_last_block(nir_function_impl * impl)2993 nir_impl_last_block(nir_function_impl *impl)
2994 {
2995 return (nir_block *) impl->body.tail_sentinel.prev;
2996 }
2997
2998 static inline nir_cf_node *
nir_cf_node_next(nir_cf_node * node)2999 nir_cf_node_next(nir_cf_node *node)
3000 {
3001 struct exec_node *next = exec_node_get_next(&node->node);
3002 if (exec_node_is_tail_sentinel(next))
3003 return NULL;
3004 else
3005 return exec_node_data(nir_cf_node, next, node);
3006 }
3007
3008 static inline nir_cf_node *
nir_cf_node_prev(nir_cf_node * node)3009 nir_cf_node_prev(nir_cf_node *node)
3010 {
3011 struct exec_node *prev = exec_node_get_prev(&node->node);
3012 if (exec_node_is_head_sentinel(prev))
3013 return NULL;
3014 else
3015 return exec_node_data(nir_cf_node, prev, node);
3016 }
3017
3018 static inline bool
nir_cf_node_is_first(const nir_cf_node * node)3019 nir_cf_node_is_first(const nir_cf_node *node)
3020 {
3021 return exec_node_is_head_sentinel(node->node.prev);
3022 }
3023
3024 static inline bool
nir_cf_node_is_last(const nir_cf_node * node)3025 nir_cf_node_is_last(const nir_cf_node *node)
3026 {
3027 return exec_node_is_tail_sentinel(node->node.next);
3028 }
3029
NIR_DEFINE_CAST(nir_cf_node_as_block,nir_cf_node,nir_block,cf_node,type,nir_cf_node_block)3030 NIR_DEFINE_CAST(nir_cf_node_as_block, nir_cf_node, nir_block, cf_node,
3031 type, nir_cf_node_block)
3032 NIR_DEFINE_CAST(nir_cf_node_as_if, nir_cf_node, nir_if, cf_node,
3033 type, nir_cf_node_if)
3034 NIR_DEFINE_CAST(nir_cf_node_as_loop, nir_cf_node, nir_loop, cf_node,
3035 type, nir_cf_node_loop)
3036 NIR_DEFINE_CAST(nir_cf_node_as_function, nir_cf_node,
3037 nir_function_impl, cf_node, type, nir_cf_node_function)
3038
3039 static inline nir_block *
3040 nir_if_first_then_block(nir_if *if_stmt)
3041 {
3042 struct exec_node *head = exec_list_get_head(&if_stmt->then_list);
3043 return nir_cf_node_as_block(exec_node_data(nir_cf_node, head, node));
3044 }
3045
3046 static inline nir_block *
nir_if_last_then_block(nir_if * if_stmt)3047 nir_if_last_then_block(nir_if *if_stmt)
3048 {
3049 struct exec_node *tail = exec_list_get_tail(&if_stmt->then_list);
3050 return nir_cf_node_as_block(exec_node_data(nir_cf_node, tail, node));
3051 }
3052
3053 static inline nir_block *
nir_if_first_else_block(nir_if * if_stmt)3054 nir_if_first_else_block(nir_if *if_stmt)
3055 {
3056 struct exec_node *head = exec_list_get_head(&if_stmt->else_list);
3057 return nir_cf_node_as_block(exec_node_data(nir_cf_node, head, node));
3058 }
3059
3060 static inline nir_block *
nir_if_last_else_block(nir_if * if_stmt)3061 nir_if_last_else_block(nir_if *if_stmt)
3062 {
3063 struct exec_node *tail = exec_list_get_tail(&if_stmt->else_list);
3064 return nir_cf_node_as_block(exec_node_data(nir_cf_node, tail, node));
3065 }
3066
3067 static inline nir_block *
nir_loop_first_block(nir_loop * loop)3068 nir_loop_first_block(nir_loop *loop)
3069 {
3070 struct exec_node *head = exec_list_get_head(&loop->body);
3071 return nir_cf_node_as_block(exec_node_data(nir_cf_node, head, node));
3072 }
3073
3074 static inline nir_block *
nir_loop_last_block(nir_loop * loop)3075 nir_loop_last_block(nir_loop *loop)
3076 {
3077 struct exec_node *tail = exec_list_get_tail(&loop->body);
3078 return nir_cf_node_as_block(exec_node_data(nir_cf_node, tail, node));
3079 }
3080
3081 /**
3082 * Return true if this list of cf_nodes contains a single empty block.
3083 */
3084 static inline bool
nir_cf_list_is_empty_block(struct exec_list * cf_list)3085 nir_cf_list_is_empty_block(struct exec_list *cf_list)
3086 {
3087 if (exec_list_is_singular(cf_list)) {
3088 struct exec_node *head = exec_list_get_head(cf_list);
3089 nir_block *block =
3090 nir_cf_node_as_block(exec_node_data(nir_cf_node, head, node));
3091 return exec_list_is_empty(&block->instr_list);
3092 }
3093 return false;
3094 }
3095
3096 typedef struct {
3097 uint8_t num_components;
3098 uint8_t bit_size;
3099 } nir_parameter;
3100
3101 typedef struct nir_printf_info {
3102 unsigned num_args;
3103 unsigned *arg_sizes;
3104 unsigned string_size;
3105 char *strings;
3106 } nir_printf_info;
3107
3108 typedef struct nir_function {
3109 struct exec_node node;
3110
3111 const char *name;
3112 struct nir_shader *shader;
3113
3114 unsigned num_params;
3115 nir_parameter *params;
3116
3117 /** The implementation of this function.
3118 *
3119 * If the function is only declared and not implemented, this is NULL.
3120 */
3121 nir_function_impl *impl;
3122
3123 bool is_entrypoint;
3124 bool is_preamble;
3125 } nir_function;
3126
3127 typedef enum {
3128 nir_lower_imul64 = (1 << 0),
3129 nir_lower_isign64 = (1 << 1),
3130 /** Lower all int64 modulus and division opcodes */
3131 nir_lower_divmod64 = (1 << 2),
3132 /** Lower all 64-bit umul_high and imul_high opcodes */
3133 nir_lower_imul_high64 = (1 << 3),
3134 nir_lower_mov64 = (1 << 4),
3135 nir_lower_icmp64 = (1 << 5),
3136 nir_lower_iadd64 = (1 << 6),
3137 nir_lower_iabs64 = (1 << 7),
3138 nir_lower_ineg64 = (1 << 8),
3139 nir_lower_logic64 = (1 << 9),
3140 nir_lower_minmax64 = (1 << 10),
3141 nir_lower_shift64 = (1 << 11),
3142 nir_lower_imul_2x32_64 = (1 << 12),
3143 nir_lower_extract64 = (1 << 13),
3144 nir_lower_ufind_msb64 = (1 << 14),
3145 nir_lower_bit_count64 = (1 << 15),
3146 nir_lower_subgroup_shuffle64 = (1 << 16),
3147 nir_lower_scan_reduce_bitwise64 = (1 << 17),
3148 nir_lower_scan_reduce_iadd64 = (1 << 18),
3149 nir_lower_vote_ieq64 = (1 << 19),
3150 nir_lower_usub_sat64 = (1 << 20),
3151 nir_lower_iadd_sat64 = (1 << 21),
3152 } nir_lower_int64_options;
3153
3154 typedef enum {
3155 nir_lower_drcp = (1 << 0),
3156 nir_lower_dsqrt = (1 << 1),
3157 nir_lower_drsq = (1 << 2),
3158 nir_lower_dtrunc = (1 << 3),
3159 nir_lower_dfloor = (1 << 4),
3160 nir_lower_dceil = (1 << 5),
3161 nir_lower_dfract = (1 << 6),
3162 nir_lower_dround_even = (1 << 7),
3163 nir_lower_dmod = (1 << 8),
3164 nir_lower_dsub = (1 << 9),
3165 nir_lower_ddiv = (1 << 10),
3166 nir_lower_fp64_full_software = (1 << 11),
3167 } nir_lower_doubles_options;
3168
3169 typedef enum {
3170 nir_divergence_single_prim_per_subgroup = (1 << 0),
3171 nir_divergence_single_patch_per_tcs_subgroup = (1 << 1),
3172 nir_divergence_single_patch_per_tes_subgroup = (1 << 2),
3173 nir_divergence_view_index_uniform = (1 << 3),
3174 nir_divergence_single_frag_shading_rate_per_subgroup = (1 << 4),
3175 nir_divergence_multiple_workgroup_per_compute_subgroup = (1 << 5),
3176 } nir_divergence_options;
3177
3178 typedef enum {
3179 nir_pack_varying_interp_mode_none = (1 << 0),
3180 nir_pack_varying_interp_mode_smooth = (1 << 1),
3181 nir_pack_varying_interp_mode_flat = (1 << 2),
3182 nir_pack_varying_interp_mode_noperspective = (1 << 3),
3183 nir_pack_varying_interp_loc_sample = (1 << 16),
3184 nir_pack_varying_interp_loc_centroid = (1 << 17),
3185 nir_pack_varying_interp_loc_center = (1 << 18),
3186 } nir_pack_varying_options;
3187
3188 /** An instruction filtering callback
3189 *
3190 * Returns true if the instruction should be processed and false otherwise.
3191 */
3192 typedef bool (*nir_instr_filter_cb)(const nir_instr *, const void *);
3193
3194 typedef struct nir_shader_compiler_options {
3195 bool lower_fdiv;
3196 bool lower_ffma16;
3197 bool lower_ffma32;
3198 bool lower_ffma64;
3199 bool fuse_ffma16;
3200 bool fuse_ffma32;
3201 bool fuse_ffma64;
3202 bool lower_flrp16;
3203 bool lower_flrp32;
3204 /** Lowers flrp when it does not support doubles */
3205 bool lower_flrp64;
3206 bool lower_fpow;
3207 bool lower_fsat;
3208 bool lower_fsqrt;
3209 bool lower_sincos;
3210 bool lower_fmod;
3211 /** Lowers ibitfield_extract/ubitfield_extract to ibfe/ubfe. */
3212 bool lower_bitfield_extract;
3213 /** Lowers ibitfield_extract/ubitfield_extract to compares, shifts. */
3214 bool lower_bitfield_extract_to_shifts;
3215 /** Lowers bitfield_insert to bfi/bfm */
3216 bool lower_bitfield_insert;
3217 /** Lowers bitfield_insert to compares, and shifts. */
3218 bool lower_bitfield_insert_to_shifts;
3219 /** Lowers bitfield_insert to bfm/bitfield_select. */
3220 bool lower_bitfield_insert_to_bitfield_select;
3221 /** Lowers bitfield_reverse to shifts. */
3222 bool lower_bitfield_reverse;
3223 /** Lowers bit_count to shifts. */
3224 bool lower_bit_count;
3225 /** Lowers ifind_msb to compare and ufind_msb */
3226 bool lower_ifind_msb;
3227 /** Lowers ifind_msb and ufind_msb to reverse variants */
3228 bool lower_find_msb_to_reverse;
3229 /** Lowers find_lsb to ufind_msb and logic ops */
3230 bool lower_find_lsb;
3231 bool lower_uadd_carry;
3232 bool lower_usub_borrow;
3233 /** Lowers imul_high/umul_high to 16-bit multiplies and carry operations. */
3234 bool lower_mul_high;
3235 /** lowers fneg to fmul(x, -1.0). Driver must call nir_opt_algebraic_late() */
3236 bool lower_fneg;
3237 /** lowers ineg to isub. Driver must call nir_opt_algebraic_late(). */
3238 bool lower_ineg;
3239 /** lowers fisnormal to alu ops. */
3240 bool lower_fisnormal;
3241
3242 /* lower {slt,sge,seq,sne} to {flt,fge,feq,fneu} + b2f: */
3243 bool lower_scmp;
3244
3245 /* lower b/fall_equalN/b/fany_nequalN (ex:fany_nequal4 to sne+fdot4+fsat) */
3246 bool lower_vector_cmp;
3247
3248 /** enable rules to avoid bit ops */
3249 bool lower_bitops;
3250
3251 /** enables rules to lower isign to imin+imax */
3252 bool lower_isign;
3253
3254 /** enables rules to lower fsign to fsub and flt */
3255 bool lower_fsign;
3256
3257 /** enables rules to lower iabs to ineg+imax */
3258 bool lower_iabs;
3259
3260 /** enable rules that avoid generating umax from signed integer ops */
3261 bool lower_umax;
3262
3263 /** enable rules that avoid generating umin from signed integer ops */
3264 bool lower_umin;
3265
3266 /* lower fdph to fdot4 */
3267 bool lower_fdph;
3268
3269 /** lower fdot to fmul and fsum/fadd. */
3270 bool lower_fdot;
3271
3272 /* Does the native fdot instruction replicate its result for four
3273 * components? If so, then opt_algebraic_late will turn all fdotN
3274 * instructions into fdotN_replicated instructions.
3275 */
3276 bool fdot_replicates;
3277
3278 /** lowers ffloor to fsub+ffract: */
3279 bool lower_ffloor;
3280
3281 /** lowers ffract to fsub+ffloor: */
3282 bool lower_ffract;
3283
3284 /** lowers fceil to fneg+ffloor+fneg: */
3285 bool lower_fceil;
3286
3287 bool lower_ftrunc;
3288
3289 bool lower_ldexp;
3290
3291 bool lower_pack_half_2x16;
3292 bool lower_pack_unorm_2x16;
3293 bool lower_pack_snorm_2x16;
3294 bool lower_pack_unorm_4x8;
3295 bool lower_pack_snorm_4x8;
3296 bool lower_pack_64_2x32;
3297 bool lower_pack_64_4x16;
3298 bool lower_pack_32_2x16;
3299 bool lower_pack_64_2x32_split;
3300 bool lower_pack_32_2x16_split;
3301 bool lower_unpack_half_2x16;
3302 bool lower_unpack_unorm_2x16;
3303 bool lower_unpack_snorm_2x16;
3304 bool lower_unpack_unorm_4x8;
3305 bool lower_unpack_snorm_4x8;
3306 bool lower_unpack_64_2x32_split;
3307 bool lower_unpack_32_2x16_split;
3308
3309 bool lower_pack_split;
3310
3311 bool lower_extract_byte;
3312 bool lower_extract_word;
3313 bool lower_insert_byte;
3314 bool lower_insert_word;
3315
3316 bool lower_all_io_to_temps;
3317 bool lower_all_io_to_elements;
3318
3319 /* Indicates that the driver only has zero-based vertex id */
3320 bool vertex_id_zero_based;
3321
3322 /**
3323 * If enabled, gl_BaseVertex will be lowered as:
3324 * is_indexed_draw (~0/0) & firstvertex
3325 */
3326 bool lower_base_vertex;
3327
3328 /**
3329 * If enabled, gl_HelperInvocation will be lowered as:
3330 *
3331 * !((1 << sample_id) & sample_mask_in))
3332 *
3333 * This depends on some possibly hw implementation details, which may
3334 * not be true for all hw. In particular that the FS is only executed
3335 * for covered samples or for helper invocations. So, do not blindly
3336 * enable this option.
3337 *
3338 * Note: See also issue #22 in ARB_shader_image_load_store
3339 */
3340 bool lower_helper_invocation;
3341
3342 /**
3343 * Convert gl_SampleMaskIn to gl_HelperInvocation as follows:
3344 *
3345 * gl_SampleMaskIn == 0 ---> gl_HelperInvocation
3346 * gl_SampleMaskIn != 0 ---> !gl_HelperInvocation
3347 */
3348 bool optimize_sample_mask_in;
3349
3350 bool lower_cs_local_index_to_id;
3351 bool lower_cs_local_id_to_index;
3352
3353 /* Prevents lowering global_invocation_id to be in terms of workgroup_id */
3354 bool has_cs_global_id;
3355
3356 bool lower_device_index_to_zero;
3357
3358 /* Set if nir_lower_pntc_ytransform() should invert gl_PointCoord.
3359 * Either when frame buffer is flipped or GL_POINT_SPRITE_COORD_ORIGIN
3360 * is GL_LOWER_LEFT.
3361 */
3362 bool lower_wpos_pntc;
3363
3364 /**
3365 * Set if nir_op_[iu]hadd and nir_op_[iu]rhadd instructions should be
3366 * lowered to simple arithmetic.
3367 *
3368 * If this flag is set, the lowering will be applied to all bit-sizes of
3369 * these instructions.
3370 *
3371 * \sa ::lower_hadd64
3372 */
3373 bool lower_hadd;
3374
3375 /**
3376 * Set if only 64-bit nir_op_[iu]hadd and nir_op_[iu]rhadd instructions
3377 * should be lowered to simple arithmetic.
3378 *
3379 * If this flag is set, the lowering will be applied to only 64-bit
3380 * versions of these instructions.
3381 *
3382 * \sa ::lower_hadd
3383 */
3384 bool lower_hadd64;
3385
3386 /**
3387 * Set if nir_op_uadd_sat and nir_op_usub_sat should be lowered to simple
3388 * arithmetic.
3389 *
3390 * If this flag is set, the lowering will be applied to all bit-sizes of
3391 * these instructions.
3392 */
3393 bool lower_uadd_sat;
3394
3395 /**
3396 * Set if nir_op_iadd_sat and nir_op_isub_sat should be lowered to simple
3397 * arithmetic.
3398 *
3399 * If this flag is set, the lowering will be applied to all bit-sizes of
3400 * these instructions.
3401 */
3402 bool lower_iadd_sat;
3403
3404 /**
3405 * Should IO be re-vectorized? Some scalar ISAs still operate on vec4's
3406 * for IO purposes and would prefer loads/stores be vectorized.
3407 */
3408 bool vectorize_io;
3409 bool lower_to_scalar;
3410 nir_instr_filter_cb lower_to_scalar_filter;
3411
3412 /**
3413 * Whether nir_opt_vectorize should only create 16-bit 2D vectors.
3414 */
3415 bool vectorize_vec2_16bit;
3416
3417 /**
3418 * Should the linker unify inputs_read/outputs_written between adjacent
3419 * shader stages which are linked into a single program?
3420 */
3421 bool unify_interfaces;
3422
3423 /**
3424 * Should nir_lower_io() create load_interpolated_input intrinsics?
3425 *
3426 * If not, it generates regular load_input intrinsics and interpolation
3427 * information must be inferred from the list of input nir_variables.
3428 */
3429 bool use_interpolated_input_intrinsics;
3430
3431
3432 /**
3433 * Whether nir_lower_io() will lower interpolateAt functions to
3434 * load_interpolated_input intrinsics.
3435 *
3436 * Unlike use_interpolated_input_intrinsics this will only lower these
3437 * functions and leave input load intrinsics untouched.
3438 */
3439 bool lower_interpolate_at;
3440
3441 /* Lowers when 32x32->64 bit multiplication is not supported */
3442 bool lower_mul_2x32_64;
3443
3444 /* Lowers when rotate instruction is not supported */
3445 bool lower_rotate;
3446
3447 /** Backend supports ternary addition */
3448 bool has_iadd3;
3449
3450 /**
3451 * Backend supports imul24, and would like to use it (when possible)
3452 * for address/offset calculation. If true, driver should call
3453 * nir_lower_amul(). (If not set, amul will automatically be lowered
3454 * to imul.)
3455 */
3456 bool has_imul24;
3457
3458 /** Backend supports umul24, if not set umul24 will automatically be lowered
3459 * to imul with masked inputs */
3460 bool has_umul24;
3461
3462 /** Backend supports umad24, if not set umad24 will automatically be lowered
3463 * to imul with masked inputs and iadd */
3464 bool has_umad24;
3465
3466 /* Backend supports fused comapre against zero and csel */
3467 bool has_fused_comp_and_csel;
3468
3469 /** Backend supports fsub, if not set fsub will automatically be lowered to
3470 * fadd(x, fneg(y)). If true, driver should call nir_opt_algebraic_late(). */
3471 bool has_fsub;
3472
3473 /** Backend supports isub, if not set isub will automatically be lowered to
3474 * iadd(x, ineg(y)). If true, driver should call nir_opt_algebraic_late(). */
3475 bool has_isub;
3476
3477 /** Backend supports pack_32_4x8 or pack_32_4x8_split. */
3478 bool has_pack_32_4x8;
3479
3480 /** Backend supports txs, if not nir_lower_tex(..) uses txs-free variants
3481 * for rect texture lowering. */
3482 bool has_txs;
3483
3484 /** Backend supports sdot_4x8 opcodes. */
3485 bool has_sdot_4x8;
3486
3487 /** Backend supports udot_4x8 opcodes. */
3488 bool has_udot_4x8;
3489
3490 /** Backend supports sudot_4x8 opcodes. */
3491 bool has_sudot_4x8;
3492
3493 /** Backend supports sdot_2x16 and udot_2x16 opcodes. */
3494 bool has_dot_2x16;
3495
3496 /* Whether to generate only scoped_barrier intrinsics instead of the set of
3497 * memory and control barrier intrinsics based on GLSL.
3498 */
3499 bool use_scoped_barrier;
3500
3501 /** Backend supports fmulz (and ffmaz if lower_ffma32=false) */
3502 bool has_fmulz;
3503
3504 /**
3505 * Is this the Intel vec4 backend?
3506 *
3507 * Used to inhibit algebraic optimizations that are known to be harmful on
3508 * the Intel vec4 backend. This is generally applicable to any
3509 * optimization that might cause more immediate values to be used in
3510 * 3-source (e.g., ffma and flrp) instructions.
3511 */
3512 bool intel_vec4;
3513
3514 /**
3515 * For most Intel GPUs, all ternary operations such as FMA and BFE cannot
3516 * have immediates, so two to three instructions may eventually be needed.
3517 */
3518 bool avoid_ternary_with_two_constants;
3519
3520 /** Whether 8-bit ALU is supported. */
3521 bool support_8bit_alu;
3522
3523 /** Whether 16-bit ALU is supported. */
3524 bool support_16bit_alu;
3525
3526 unsigned max_unroll_iterations;
3527 unsigned max_unroll_iterations_aggressive;
3528
3529 bool lower_uniforms_to_ubo;
3530
3531 /* If the precision is ignored, backends that don't handle
3532 * different precisions when passing data between stages and use
3533 * vectorized IO can pack more varyings when linking. */
3534 bool linker_ignore_precision;
3535
3536 /**
3537 * Specifies which type of indirectly accessed variables should force
3538 * loop unrolling.
3539 */
3540 nir_variable_mode force_indirect_unrolling;
3541
3542 nir_lower_int64_options lower_int64_options;
3543 nir_lower_doubles_options lower_doubles_options;
3544 nir_divergence_options divergence_analysis_options;
3545
3546 /**
3547 * Support pack varyings with different interpolation location
3548 * (center, centroid, sample) and mode (flat, noperspective, smooth)
3549 * into same slot.
3550 */
3551 nir_pack_varying_options pack_varying_options;
3552
3553 /**
3554 * Lower load_deref/store_deref of inputs and outputs into
3555 * load_input/store_input intrinsics. This is used by nir_lower_io_passes.
3556 */
3557 bool lower_io_variables;
3558
3559 /**
3560 * Lower color inputs to load_colorN that are kind of like system values
3561 * if lower_io_variables is also set. shader_info will contain
3562 * the interpolation settings. This is used by nir_lower_io_passes.
3563 */
3564 bool lower_fs_color_inputs;
3565
3566 /**
3567 * The masks of shader stages that support indirect indexing with
3568 * load_input and store_output intrinsics. It's used when
3569 * lower_io_variables is true. This is used by nir_lower_io_passes.
3570 */
3571 uint8_t support_indirect_inputs;
3572 uint8_t support_indirect_outputs;
3573
3574 /**
3575 * Remove varying loaded from uniform, let fragment shader load the
3576 * uniform directly. GPU passing varying by memory can benifit from it
3577 * for sure; but GPU passing varying by on chip resource may not.
3578 * Because it saves on chip resource but may increase memory pressure when
3579 * fragment task is far more than vertex one, so better left it disabled.
3580 */
3581 bool lower_varying_from_uniform;
3582 } nir_shader_compiler_options;
3583
3584 typedef struct nir_shader {
3585 /** list of uniforms (nir_variable) */
3586 struct exec_list variables;
3587
3588 /** Set of driver-specific options for the shader.
3589 *
3590 * The memory for the options is expected to be kept in a single static
3591 * copy by the driver.
3592 */
3593 const struct nir_shader_compiler_options *options;
3594
3595 /** Various bits of compile-time information about a given shader */
3596 struct shader_info info;
3597
3598 struct exec_list functions; /** < list of nir_function */
3599
3600 struct list_head gc_list; /** < list of all nir_instrs allocated on the shader but not yet freed. */
3601
3602 /**
3603 * The size of the variable space for load_input_*, load_uniform_*, etc.
3604 * intrinsics. This is in back-end specific units which is likely one of
3605 * bytes, dwords, or vec4s depending on context and back-end.
3606 */
3607 unsigned num_inputs, num_uniforms, num_outputs;
3608
3609 /** Size in bytes of required implicitly bound global memory */
3610 unsigned global_mem_size;
3611
3612 /** Size in bytes of required scratch space */
3613 unsigned scratch_size;
3614
3615 /** Constant data associated with this shader.
3616 *
3617 * Constant data is loaded through load_constant intrinsics (as compared to
3618 * the NIR load_const instructions which have the constant value inlined
3619 * into them). This is usually generated by nir_opt_large_constants (so
3620 * shaders don't have to load_const into a temporary array when they want
3621 * to indirect on a const array).
3622 */
3623 void *constant_data;
3624 /** Size of the constant data associated with the shader, in bytes */
3625 unsigned constant_data_size;
3626
3627 unsigned printf_info_count;
3628 nir_printf_info *printf_info;
3629 } nir_shader;
3630
3631 #define nir_foreach_function(func, shader) \
3632 foreach_list_typed(nir_function, func, node, &(shader)->functions)
3633
3634 static inline nir_function_impl *
nir_shader_get_entrypoint(const nir_shader * shader)3635 nir_shader_get_entrypoint(const nir_shader *shader)
3636 {
3637 nir_function *func = NULL;
3638
3639 nir_foreach_function(function, shader) {
3640 assert(func == NULL);
3641 if (function->is_entrypoint) {
3642 func = function;
3643 #ifndef NDEBUG
3644 break;
3645 #endif
3646 }
3647 }
3648
3649 if (!func)
3650 return NULL;
3651
3652 assert(func->num_params == 0);
3653 assert(func->impl);
3654 return func->impl;
3655 }
3656
3657 nir_shader *nir_shader_create(void *mem_ctx,
3658 gl_shader_stage stage,
3659 const nir_shader_compiler_options *options,
3660 shader_info *si);
3661
3662 nir_register *nir_local_reg_create(nir_function_impl *impl);
3663
3664 void nir_reg_remove(nir_register *reg);
3665
3666 /** Adds a variable to the appropriate list in nir_shader */
3667 void nir_shader_add_variable(nir_shader *shader, nir_variable *var);
3668
3669 static inline void
nir_function_impl_add_variable(nir_function_impl * impl,nir_variable * var)3670 nir_function_impl_add_variable(nir_function_impl *impl, nir_variable *var)
3671 {
3672 assert(var->data.mode == nir_var_function_temp);
3673 exec_list_push_tail(&impl->locals, &var->node);
3674 }
3675
3676 /** creates a variable, sets a few defaults, and adds it to the list */
3677 nir_variable *nir_variable_create(nir_shader *shader,
3678 nir_variable_mode mode,
3679 const struct glsl_type *type,
3680 const char *name);
3681 /** creates a local variable and adds it to the list */
3682 nir_variable *nir_local_variable_create(nir_function_impl *impl,
3683 const struct glsl_type *type,
3684 const char *name);
3685
3686 nir_variable *nir_find_variable_with_location(nir_shader *shader,
3687 nir_variable_mode mode,
3688 unsigned location);
3689
3690 nir_variable *nir_find_variable_with_driver_location(nir_shader *shader,
3691 nir_variable_mode mode,
3692 unsigned location);
3693
3694 void nir_sort_variables_with_modes(nir_shader *shader,
3695 int (*compar)(const nir_variable *,
3696 const nir_variable *),
3697 nir_variable_mode modes);
3698
3699 /** creates a function and adds it to the shader's list of functions */
3700 nir_function *nir_function_create(nir_shader *shader, const char *name);
3701
3702 nir_function_impl *nir_function_impl_create(nir_function *func);
3703 /** creates a function_impl that isn't tied to any particular function */
3704 nir_function_impl *nir_function_impl_create_bare(nir_shader *shader);
3705
3706 nir_block *nir_block_create(nir_shader *shader);
3707 nir_if *nir_if_create(nir_shader *shader);
3708 nir_loop *nir_loop_create(nir_shader *shader);
3709
3710 nir_function_impl *nir_cf_node_get_function(nir_cf_node *node);
3711
3712 /** requests that the given pieces of metadata be generated */
3713 void nir_metadata_require(nir_function_impl *impl, nir_metadata required, ...);
3714 /** dirties all but the preserved metadata */
3715 void nir_metadata_preserve(nir_function_impl *impl, nir_metadata preserved);
3716 /** Preserves all metadata for the given shader */
3717 void nir_shader_preserve_all_metadata(nir_shader *shader);
3718
3719 /** creates an instruction with default swizzle/writemask/etc. with NULL registers */
3720 nir_alu_instr *nir_alu_instr_create(nir_shader *shader, nir_op op);
3721
3722 nir_deref_instr *nir_deref_instr_create(nir_shader *shader,
3723 nir_deref_type deref_type);
3724
3725 nir_jump_instr *nir_jump_instr_create(nir_shader *shader, nir_jump_type type);
3726
3727 nir_load_const_instr *nir_load_const_instr_create(nir_shader *shader,
3728 unsigned num_components,
3729 unsigned bit_size);
3730
3731 nir_intrinsic_instr *nir_intrinsic_instr_create(nir_shader *shader,
3732 nir_intrinsic_op op);
3733
3734 nir_call_instr *nir_call_instr_create(nir_shader *shader,
3735 nir_function *callee);
3736
3737 /** Creates a NIR texture instruction */
3738 nir_tex_instr *nir_tex_instr_create(nir_shader *shader, unsigned num_srcs);
3739
3740 nir_phi_instr *nir_phi_instr_create(nir_shader *shader);
3741 nir_phi_src *nir_phi_instr_add_src(nir_phi_instr *instr, nir_block *pred, nir_src src);
3742
3743 nir_parallel_copy_instr *nir_parallel_copy_instr_create(nir_shader *shader);
3744
3745 nir_ssa_undef_instr *nir_ssa_undef_instr_create(nir_shader *shader,
3746 unsigned num_components,
3747 unsigned bit_size);
3748
3749 nir_const_value nir_alu_binop_identity(nir_op binop, unsigned bit_size);
3750
3751 /**
3752 * NIR Cursors and Instruction Insertion API
3753 * @{
3754 *
3755 * A tiny struct representing a point to insert/extract instructions or
3756 * control flow nodes. Helps reduce the combinatorial explosion of possible
3757 * points to insert/extract.
3758 *
3759 * \sa nir_control_flow.h
3760 */
3761 typedef enum {
3762 nir_cursor_before_block,
3763 nir_cursor_after_block,
3764 nir_cursor_before_instr,
3765 nir_cursor_after_instr,
3766 } nir_cursor_option;
3767
3768 typedef struct {
3769 nir_cursor_option option;
3770 union {
3771 nir_block *block;
3772 nir_instr *instr;
3773 };
3774 } nir_cursor;
3775
3776 static inline nir_block *
nir_cursor_current_block(nir_cursor cursor)3777 nir_cursor_current_block(nir_cursor cursor)
3778 {
3779 if (cursor.option == nir_cursor_before_instr ||
3780 cursor.option == nir_cursor_after_instr) {
3781 return cursor.instr->block;
3782 } else {
3783 return cursor.block;
3784 }
3785 }
3786
3787 bool nir_cursors_equal(nir_cursor a, nir_cursor b);
3788
3789 static inline nir_cursor
nir_before_block(nir_block * block)3790 nir_before_block(nir_block *block)
3791 {
3792 nir_cursor cursor;
3793 cursor.option = nir_cursor_before_block;
3794 cursor.block = block;
3795 return cursor;
3796 }
3797
3798 static inline nir_cursor
nir_after_block(nir_block * block)3799 nir_after_block(nir_block *block)
3800 {
3801 nir_cursor cursor;
3802 cursor.option = nir_cursor_after_block;
3803 cursor.block = block;
3804 return cursor;
3805 }
3806
3807 static inline nir_cursor
nir_before_instr(nir_instr * instr)3808 nir_before_instr(nir_instr *instr)
3809 {
3810 nir_cursor cursor;
3811 cursor.option = nir_cursor_before_instr;
3812 cursor.instr = instr;
3813 return cursor;
3814 }
3815
3816 static inline nir_cursor
nir_after_instr(nir_instr * instr)3817 nir_after_instr(nir_instr *instr)
3818 {
3819 nir_cursor cursor;
3820 cursor.option = nir_cursor_after_instr;
3821 cursor.instr = instr;
3822 return cursor;
3823 }
3824
3825 static inline nir_cursor
nir_before_block_after_phis(nir_block * block)3826 nir_before_block_after_phis(nir_block *block)
3827 {
3828 nir_phi_instr *last_phi = nir_block_last_phi_instr(block);
3829 if (last_phi)
3830 return nir_after_instr(&last_phi->instr);
3831 else
3832 return nir_before_block(block);
3833 }
3834
3835 static inline nir_cursor
nir_after_block_before_jump(nir_block * block)3836 nir_after_block_before_jump(nir_block *block)
3837 {
3838 nir_instr *last_instr = nir_block_last_instr(block);
3839 if (last_instr && last_instr->type == nir_instr_type_jump) {
3840 return nir_before_instr(last_instr);
3841 } else {
3842 return nir_after_block(block);
3843 }
3844 }
3845
3846 static inline nir_cursor
nir_before_src(nir_src * src,bool is_if_condition)3847 nir_before_src(nir_src *src, bool is_if_condition)
3848 {
3849 if (is_if_condition) {
3850 nir_block *prev_block =
3851 nir_cf_node_as_block(nir_cf_node_prev(&src->parent_if->cf_node));
3852 assert(!nir_block_ends_in_jump(prev_block));
3853 return nir_after_block(prev_block);
3854 } else if (src->parent_instr->type == nir_instr_type_phi) {
3855 #ifndef NDEBUG
3856 nir_phi_instr *cond_phi = nir_instr_as_phi(src->parent_instr);
3857 bool found = false;
3858 nir_foreach_phi_src(phi_src, cond_phi) {
3859 if (phi_src->src.ssa == src->ssa) {
3860 found = true;
3861 break;
3862 }
3863 }
3864 assert(found);
3865 #endif
3866 /* The LIST_ENTRY macro is a generic container-of macro, it just happens
3867 * to have a more specific name.
3868 */
3869 nir_phi_src *phi_src = LIST_ENTRY(nir_phi_src, src, src);
3870 return nir_after_block_before_jump(phi_src->pred);
3871 } else {
3872 return nir_before_instr(src->parent_instr);
3873 }
3874 }
3875
3876 static inline nir_cursor
nir_before_cf_node(nir_cf_node * node)3877 nir_before_cf_node(nir_cf_node *node)
3878 {
3879 if (node->type == nir_cf_node_block)
3880 return nir_before_block(nir_cf_node_as_block(node));
3881
3882 return nir_after_block(nir_cf_node_as_block(nir_cf_node_prev(node)));
3883 }
3884
3885 static inline nir_cursor
nir_after_cf_node(nir_cf_node * node)3886 nir_after_cf_node(nir_cf_node *node)
3887 {
3888 if (node->type == nir_cf_node_block)
3889 return nir_after_block(nir_cf_node_as_block(node));
3890
3891 return nir_before_block(nir_cf_node_as_block(nir_cf_node_next(node)));
3892 }
3893
3894 static inline nir_cursor
nir_after_phis(nir_block * block)3895 nir_after_phis(nir_block *block)
3896 {
3897 nir_foreach_instr(instr, block) {
3898 if (instr->type != nir_instr_type_phi)
3899 return nir_before_instr(instr);
3900 }
3901 return nir_after_block(block);
3902 }
3903
3904 static inline nir_cursor
nir_after_instr_and_phis(nir_instr * instr)3905 nir_after_instr_and_phis(nir_instr *instr)
3906 {
3907 if (instr->type == nir_instr_type_phi)
3908 return nir_after_phis(instr->block);
3909 else
3910 return nir_after_instr(instr);
3911 }
3912
3913 static inline nir_cursor
nir_after_cf_node_and_phis(nir_cf_node * node)3914 nir_after_cf_node_and_phis(nir_cf_node *node)
3915 {
3916 if (node->type == nir_cf_node_block)
3917 return nir_after_block(nir_cf_node_as_block(node));
3918
3919 nir_block *block = nir_cf_node_as_block(nir_cf_node_next(node));
3920
3921 return nir_after_phis(block);
3922 }
3923
3924 static inline nir_cursor
nir_before_cf_list(struct exec_list * cf_list)3925 nir_before_cf_list(struct exec_list *cf_list)
3926 {
3927 nir_cf_node *first_node = exec_node_data(nir_cf_node,
3928 exec_list_get_head(cf_list), node);
3929 return nir_before_cf_node(first_node);
3930 }
3931
3932 static inline nir_cursor
nir_after_cf_list(struct exec_list * cf_list)3933 nir_after_cf_list(struct exec_list *cf_list)
3934 {
3935 nir_cf_node *last_node = exec_node_data(nir_cf_node,
3936 exec_list_get_tail(cf_list), node);
3937 return nir_after_cf_node(last_node);
3938 }
3939
3940 /**
3941 * Insert a NIR instruction at the given cursor.
3942 *
3943 * Note: This does not update the cursor.
3944 */
3945 void nir_instr_insert(nir_cursor cursor, nir_instr *instr);
3946
3947 bool nir_instr_move(nir_cursor cursor, nir_instr *instr);
3948
3949 static inline void
nir_instr_insert_before(nir_instr * instr,nir_instr * before)3950 nir_instr_insert_before(nir_instr *instr, nir_instr *before)
3951 {
3952 nir_instr_insert(nir_before_instr(instr), before);
3953 }
3954
3955 static inline void
nir_instr_insert_after(nir_instr * instr,nir_instr * after)3956 nir_instr_insert_after(nir_instr *instr, nir_instr *after)
3957 {
3958 nir_instr_insert(nir_after_instr(instr), after);
3959 }
3960
3961 static inline void
nir_instr_insert_before_block(nir_block * block,nir_instr * before)3962 nir_instr_insert_before_block(nir_block *block, nir_instr *before)
3963 {
3964 nir_instr_insert(nir_before_block(block), before);
3965 }
3966
3967 static inline void
nir_instr_insert_after_block(nir_block * block,nir_instr * after)3968 nir_instr_insert_after_block(nir_block *block, nir_instr *after)
3969 {
3970 nir_instr_insert(nir_after_block(block), after);
3971 }
3972
3973 static inline void
nir_instr_insert_before_cf(nir_cf_node * node,nir_instr * before)3974 nir_instr_insert_before_cf(nir_cf_node *node, nir_instr *before)
3975 {
3976 nir_instr_insert(nir_before_cf_node(node), before);
3977 }
3978
3979 static inline void
nir_instr_insert_after_cf(nir_cf_node * node,nir_instr * after)3980 nir_instr_insert_after_cf(nir_cf_node *node, nir_instr *after)
3981 {
3982 nir_instr_insert(nir_after_cf_node(node), after);
3983 }
3984
3985 static inline void
nir_instr_insert_before_cf_list(struct exec_list * list,nir_instr * before)3986 nir_instr_insert_before_cf_list(struct exec_list *list, nir_instr *before)
3987 {
3988 nir_instr_insert(nir_before_cf_list(list), before);
3989 }
3990
3991 static inline void
nir_instr_insert_after_cf_list(struct exec_list * list,nir_instr * after)3992 nir_instr_insert_after_cf_list(struct exec_list *list, nir_instr *after)
3993 {
3994 nir_instr_insert(nir_after_cf_list(list), after);
3995 }
3996
3997 void nir_instr_remove_v(nir_instr *instr);
3998 void nir_instr_free(nir_instr *instr);
3999 void nir_instr_free_list(struct exec_list *list);
4000
4001 static inline nir_cursor
nir_instr_remove(nir_instr * instr)4002 nir_instr_remove(nir_instr *instr)
4003 {
4004 nir_cursor cursor;
4005 nir_instr *prev = nir_instr_prev(instr);
4006 if (prev) {
4007 cursor = nir_after_instr(prev);
4008 } else {
4009 cursor = nir_before_block(instr->block);
4010 }
4011 nir_instr_remove_v(instr);
4012 return cursor;
4013 }
4014
4015 nir_cursor nir_instr_free_and_dce(nir_instr *instr);
4016
4017 /** @} */
4018
4019 nir_ssa_def *nir_instr_ssa_def(nir_instr *instr);
4020 bool nir_instr_def_is_register(nir_instr *instr);
4021
4022 typedef bool (*nir_foreach_ssa_def_cb)(nir_ssa_def *def, void *state);
4023 typedef bool (*nir_foreach_dest_cb)(nir_dest *dest, void *state);
4024 typedef bool (*nir_foreach_src_cb)(nir_src *src, void *state);
4025 bool nir_foreach_ssa_def(nir_instr *instr, nir_foreach_ssa_def_cb cb,
4026 void *state);
4027 static inline bool nir_foreach_dest(nir_instr *instr, nir_foreach_dest_cb cb, void *state);
4028 static inline bool nir_foreach_src(nir_instr *instr, nir_foreach_src_cb cb, void *state);
4029 bool nir_foreach_phi_src_leaving_block(nir_block *instr,
4030 nir_foreach_src_cb cb,
4031 void *state);
4032
4033 nir_const_value *nir_src_as_const_value(nir_src src);
4034
4035 #define NIR_SRC_AS_(name, c_type, type_enum, cast_macro) \
4036 static inline c_type * \
4037 nir_src_as_ ## name (nir_src src) \
4038 { \
4039 return src.is_ssa && src.ssa->parent_instr->type == type_enum \
4040 ? cast_macro(src.ssa->parent_instr) : NULL; \
4041 }
4042
4043 NIR_SRC_AS_(alu_instr, nir_alu_instr, nir_instr_type_alu, nir_instr_as_alu)
4044 NIR_SRC_AS_(intrinsic, nir_intrinsic_instr,
4045 nir_instr_type_intrinsic, nir_instr_as_intrinsic)
4046 NIR_SRC_AS_(deref, nir_deref_instr, nir_instr_type_deref, nir_instr_as_deref)
4047
4048 bool nir_src_is_always_uniform(nir_src src);
4049 bool nir_srcs_equal(nir_src src1, nir_src src2);
4050 bool nir_instrs_equal(const nir_instr *instr1, const nir_instr *instr2);
4051
4052 static inline void
nir_instr_rewrite_src_ssa(ASSERTED nir_instr * instr,nir_src * src,nir_ssa_def * new_ssa)4053 nir_instr_rewrite_src_ssa(ASSERTED nir_instr *instr,
4054 nir_src *src, nir_ssa_def *new_ssa)
4055 {
4056 assert(src->parent_instr == instr);
4057 assert(src->is_ssa && src->ssa);
4058 list_del(&src->use_link);
4059 src->ssa = new_ssa;
4060 list_addtail(&src->use_link, &new_ssa->uses);
4061 }
4062
4063 void nir_instr_rewrite_src(nir_instr *instr, nir_src *src, nir_src new_src);
4064 void nir_instr_move_src(nir_instr *dest_instr, nir_src *dest, nir_src *src);
4065
4066 static inline void
nir_if_rewrite_condition_ssa(ASSERTED nir_if * if_stmt,nir_src * src,nir_ssa_def * new_ssa)4067 nir_if_rewrite_condition_ssa(ASSERTED nir_if *if_stmt,
4068 nir_src *src, nir_ssa_def *new_ssa)
4069 {
4070 assert(src->parent_if == if_stmt);
4071 assert(src->is_ssa && src->ssa);
4072 list_del(&src->use_link);
4073 src->ssa = new_ssa;
4074 list_addtail(&src->use_link, &new_ssa->if_uses);
4075 }
4076
4077 void nir_if_rewrite_condition(nir_if *if_stmt, nir_src new_src);
4078 void nir_instr_rewrite_dest(nir_instr *instr, nir_dest *dest,
4079 nir_dest new_dest);
4080
4081 void nir_ssa_dest_init(nir_instr *instr, nir_dest *dest,
4082 unsigned num_components, unsigned bit_size,
4083 const char *name);
4084 void nir_ssa_def_init(nir_instr *instr, nir_ssa_def *def,
4085 unsigned num_components, unsigned bit_size);
4086 static inline void
nir_ssa_dest_init_for_type(nir_instr * instr,nir_dest * dest,const struct glsl_type * type,const char * name)4087 nir_ssa_dest_init_for_type(nir_instr *instr, nir_dest *dest,
4088 const struct glsl_type *type,
4089 const char *name)
4090 {
4091 assert(glsl_type_is_vector_or_scalar(type));
4092 nir_ssa_dest_init(instr, dest, glsl_get_components(type),
4093 glsl_get_bit_size(type), name);
4094 }
4095 void nir_ssa_def_rewrite_uses(nir_ssa_def *def, nir_ssa_def *new_ssa);
4096 void nir_ssa_def_rewrite_uses_src(nir_ssa_def *def, nir_src new_src);
4097 void nir_ssa_def_rewrite_uses_after(nir_ssa_def *def, nir_ssa_def *new_ssa,
4098 nir_instr *after_me);
4099
4100 nir_component_mask_t nir_src_components_read(const nir_src *src);
4101 nir_component_mask_t nir_ssa_def_components_read(const nir_ssa_def *def);
4102
4103 static inline bool
nir_ssa_def_is_unused(nir_ssa_def * ssa)4104 nir_ssa_def_is_unused(nir_ssa_def *ssa)
4105 {
4106 return list_is_empty(&ssa->uses) && list_is_empty(&ssa->if_uses);
4107 }
4108
4109
4110 /** Returns the next block, disregarding structure
4111 *
4112 * The ordering is deterministic but has no guarantees beyond that. In
4113 * particular, it is not guaranteed to be dominance-preserving.
4114 */
4115 nir_block *nir_block_unstructured_next(nir_block *block);
4116 nir_block *nir_unstructured_start_block(nir_function_impl *impl);
4117
4118 #define nir_foreach_block_unstructured(block, impl) \
4119 for (nir_block *block = nir_unstructured_start_block(impl); block != NULL; \
4120 block = nir_block_unstructured_next(block))
4121
4122 #define nir_foreach_block_unstructured_safe(block, impl) \
4123 for (nir_block *block = nir_unstructured_start_block(impl), \
4124 *next = nir_block_unstructured_next(block); \
4125 block != NULL; \
4126 block = next, next = nir_block_unstructured_next(block))
4127
4128 /*
4129 * finds the next basic block in source-code order, returns NULL if there is
4130 * none
4131 */
4132
4133 nir_block *nir_block_cf_tree_next(nir_block *block);
4134
4135 /* Performs the opposite of nir_block_cf_tree_next() */
4136
4137 nir_block *nir_block_cf_tree_prev(nir_block *block);
4138
4139 /* Gets the first block in a CF node in source-code order */
4140
4141 nir_block *nir_cf_node_cf_tree_first(nir_cf_node *node);
4142
4143 /* Gets the last block in a CF node in source-code order */
4144
4145 nir_block *nir_cf_node_cf_tree_last(nir_cf_node *node);
4146
4147 /* Gets the next block after a CF node in source-code order */
4148
4149 nir_block *nir_cf_node_cf_tree_next(nir_cf_node *node);
4150
4151 /* Macros for loops that visit blocks in source-code order */
4152
4153 #define nir_foreach_block(block, impl) \
4154 for (nir_block *block = nir_start_block(impl); block != NULL; \
4155 block = nir_block_cf_tree_next(block))
4156
4157 #define nir_foreach_block_safe(block, impl) \
4158 for (nir_block *block = nir_start_block(impl), \
4159 *next = nir_block_cf_tree_next(block); \
4160 block != NULL; \
4161 block = next, next = nir_block_cf_tree_next(block))
4162
4163 #define nir_foreach_block_reverse(block, impl) \
4164 for (nir_block *block = nir_impl_last_block(impl); block != NULL; \
4165 block = nir_block_cf_tree_prev(block))
4166
4167 #define nir_foreach_block_reverse_safe(block, impl) \
4168 for (nir_block *block = nir_impl_last_block(impl), \
4169 *prev = nir_block_cf_tree_prev(block); \
4170 block != NULL; \
4171 block = prev, prev = nir_block_cf_tree_prev(block))
4172
4173 #define nir_foreach_block_in_cf_node(block, node) \
4174 for (nir_block *block = nir_cf_node_cf_tree_first(node); \
4175 block != nir_cf_node_cf_tree_next(node); \
4176 block = nir_block_cf_tree_next(block))
4177
4178 /* If the following CF node is an if, this function returns that if.
4179 * Otherwise, it returns NULL.
4180 */
4181 nir_if *nir_block_get_following_if(nir_block *block);
4182
4183 nir_loop *nir_block_get_following_loop(nir_block *block);
4184
4185 nir_block **nir_block_get_predecessors_sorted(const nir_block *block, void *mem_ctx);
4186
4187 void nir_index_local_regs(nir_function_impl *impl);
4188 void nir_index_ssa_defs(nir_function_impl *impl);
4189 unsigned nir_index_instrs(nir_function_impl *impl);
4190
4191 void nir_index_blocks(nir_function_impl *impl);
4192
4193 unsigned nir_shader_index_vars(nir_shader *shader, nir_variable_mode modes);
4194 unsigned nir_function_impl_index_vars(nir_function_impl *impl);
4195
4196 void nir_print_shader(nir_shader *shader, FILE *fp);
4197 void nir_print_shader_annotated(nir_shader *shader, FILE *fp, struct hash_table *errors);
4198 void nir_print_instr(const nir_instr *instr, FILE *fp);
4199 void nir_print_deref(const nir_deref_instr *deref, FILE *fp);
4200 void nir_log_shader_annotated_tagged(enum mesa_log_level level, const char *tag, nir_shader *shader, struct hash_table *annotations);
4201 #define nir_log_shadere(s) nir_log_shader_annotated_tagged(MESA_LOG_ERROR, (MESA_LOG_TAG), (s), NULL)
4202 #define nir_log_shaderw(s) nir_log_shader_annotated_tagged(MESA_LOG_WARN, (MESA_LOG_TAG), (s), NULL)
4203 #define nir_log_shaderi(s) nir_log_shader_annotated_tagged(MESA_LOG_INFO, (MESA_LOG_TAG), (s), NULL)
4204 #define nir_log_shader_annotated(s, annotations) nir_log_shader_annotated_tagged(MESA_LOG_ERROR, (MESA_LOG_TAG), (s), annotations)
4205
4206 char *nir_shader_as_str(nir_shader *nir, void *mem_ctx);
4207 char *nir_shader_as_str_annotated(nir_shader *nir, struct hash_table *annotations, void *mem_ctx);
4208
4209 /** Shallow clone of a single instruction. */
4210 nir_instr *nir_instr_clone(nir_shader *s, const nir_instr *orig);
4211
4212 /** Clone a single instruction, including a remap table to rewrite sources. */
4213 nir_instr *nir_instr_clone_deep(nir_shader *s, const nir_instr *orig,
4214 struct hash_table *remap_table);
4215
4216 /** Shallow clone of a single ALU instruction. */
4217 nir_alu_instr *nir_alu_instr_clone(nir_shader *s, const nir_alu_instr *orig);
4218
4219 nir_shader *nir_shader_clone(void *mem_ctx, const nir_shader *s);
4220 nir_function_impl *nir_function_impl_clone(nir_shader *shader,
4221 const nir_function_impl *fi);
4222 nir_constant *nir_constant_clone(const nir_constant *c, nir_variable *var);
4223 nir_variable *nir_variable_clone(const nir_variable *c, nir_shader *shader);
4224
4225 void nir_shader_replace(nir_shader *dest, nir_shader *src);
4226
4227 void nir_shader_serialize_deserialize(nir_shader *s);
4228
4229 #ifndef NDEBUG
4230 void nir_validate_shader(nir_shader *shader, const char *when);
4231 void nir_validate_ssa_dominance(nir_shader *shader, const char *when);
4232 void nir_metadata_set_validation_flag(nir_shader *shader);
4233 void nir_metadata_check_validation_flag(nir_shader *shader);
4234
4235 static inline bool
should_skip_nir(const char * name)4236 should_skip_nir(const char *name)
4237 {
4238 static const char *list = NULL;
4239 if (!list) {
4240 /* Comma separated list of names to skip. */
4241 list = getenv("NIR_SKIP");
4242 if (!list)
4243 list = "";
4244 }
4245
4246 if (!list[0])
4247 return false;
4248
4249 return comma_separated_list_contains(list, name);
4250 }
4251
4252 static inline bool
should_print_nir(nir_shader * shader)4253 should_print_nir(nir_shader *shader)
4254 {
4255 if (shader->info.internal ||
4256 shader->info.stage < 0 ||
4257 shader->info.stage > MESA_SHADER_KERNEL)
4258 return false;
4259
4260 return unlikely(nir_debug_print_shader[shader->info.stage]);
4261 }
4262 #else
nir_validate_shader(nir_shader * shader,const char * when)4263 static inline void nir_validate_shader(nir_shader *shader, const char *when) { (void) shader; (void)when; }
nir_validate_ssa_dominance(nir_shader * shader,const char * when)4264 static inline void nir_validate_ssa_dominance(nir_shader *shader, const char *when) { (void) shader; (void)when; }
nir_metadata_set_validation_flag(nir_shader * shader)4265 static inline void nir_metadata_set_validation_flag(nir_shader *shader) { (void) shader; }
nir_metadata_check_validation_flag(nir_shader * shader)4266 static inline void nir_metadata_check_validation_flag(nir_shader *shader) { (void) shader; }
should_skip_nir(UNUSED const char * pass_name)4267 static inline bool should_skip_nir(UNUSED const char *pass_name) { return false; }
should_print_nir(UNUSED nir_shader * shader)4268 static inline bool should_print_nir(UNUSED nir_shader *shader) { return false; }
4269 #endif /* NDEBUG */
4270
4271 #define _PASS(pass, nir, do_pass) do { \
4272 if (should_skip_nir(#pass)) { \
4273 printf("skipping %s\n", #pass); \
4274 break; \
4275 } \
4276 do_pass \
4277 if (NIR_DEBUG(CLONE)) { \
4278 nir_shader *clone = nir_shader_clone(ralloc_parent(nir), nir); \
4279 nir_shader_replace(nir, clone); \
4280 } \
4281 if (NIR_DEBUG(SERIALIZE)) { \
4282 nir_shader_serialize_deserialize(nir); \
4283 } \
4284 } while (0)
4285
4286 #define NIR_PASS(progress, nir, pass, ...) _PASS(pass, nir, \
4287 nir_metadata_set_validation_flag(nir); \
4288 if (should_print_nir(nir)) \
4289 printf("%s\n", #pass); \
4290 if (pass(nir, ##__VA_ARGS__)) { \
4291 nir_validate_shader(nir, "after " #pass); \
4292 progress = true; \
4293 if (should_print_nir(nir)) \
4294 nir_print_shader(nir, stdout); \
4295 nir_metadata_check_validation_flag(nir); \
4296 } \
4297 )
4298
4299 #define NIR_PASS_V(nir, pass, ...) _PASS(pass, nir, \
4300 if (should_print_nir(nir)) \
4301 printf("%s\n", #pass); \
4302 pass(nir, ##__VA_ARGS__); \
4303 nir_validate_shader(nir, "after " #pass); \
4304 if (should_print_nir(nir)) \
4305 nir_print_shader(nir, stdout); \
4306 )
4307
4308 #define NIR_SKIP(name) should_skip_nir(#name)
4309
4310 /** An instruction filtering callback with writemask
4311 *
4312 * Returns true if the instruction should be processed with the associated
4313 * writemask and false otherwise.
4314 */
4315 typedef bool (*nir_instr_writemask_filter_cb)(const nir_instr *,
4316 unsigned writemask, const void *);
4317
4318 /** A simple instruction lowering callback
4319 *
4320 * Many instruction lowering passes can be written as a simple function which
4321 * takes an instruction as its input and returns a sequence of instructions
4322 * that implement the consumed instruction. This function type represents
4323 * such a lowering function. When called, a function with this prototype
4324 * should either return NULL indicating that no lowering needs to be done or
4325 * emit a sequence of instructions using the provided builder (whose cursor
4326 * will already be placed after the instruction to be lowered) and return the
4327 * resulting nir_ssa_def.
4328 */
4329 typedef nir_ssa_def *(*nir_lower_instr_cb)(struct nir_builder *,
4330 nir_instr *, void *);
4331
4332 /**
4333 * Special return value for nir_lower_instr_cb when some progress occurred
4334 * (like changing an input to the instr) that didn't result in a replacement
4335 * SSA def being generated.
4336 */
4337 #define NIR_LOWER_INSTR_PROGRESS ((nir_ssa_def *)(uintptr_t)1)
4338
4339 /**
4340 * Special return value for nir_lower_instr_cb when some progress occurred
4341 * that should remove the current instruction that doesn't create an output
4342 * (like a store)
4343 */
4344
4345 #define NIR_LOWER_INSTR_PROGRESS_REPLACE ((nir_ssa_def *)(uintptr_t)2)
4346
4347 /** Iterate over all the instructions in a nir_function_impl and lower them
4348 * using the provided callbacks
4349 *
4350 * This function implements the guts of a standard lowering pass for you. It
4351 * iterates over all of the instructions in a nir_function_impl and calls the
4352 * filter callback on each one. If the filter callback returns true, it then
4353 * calls the lowering call back on the instruction. (Splitting it this way
4354 * allows us to avoid some save/restore work for instructions we know won't be
4355 * lowered.) If the instruction is dead after the lowering is complete, it
4356 * will be removed. If new instructions are added, the lowering callback will
4357 * also be called on them in case multiple lowerings are required.
4358 *
4359 * If the callback indicates that the original instruction is replaced (either
4360 * through a new SSA def or NIR_LOWER_INSTR_PROGRESS_REPLACE), then the
4361 * instruction is removed along with any now-dead SSA defs it used.
4362 *
4363 * The metadata for the nir_function_impl will also be updated. If any blocks
4364 * are added (they cannot be removed), dominance and block indices will be
4365 * invalidated.
4366 */
4367 bool nir_function_impl_lower_instructions(nir_function_impl *impl,
4368 nir_instr_filter_cb filter,
4369 nir_lower_instr_cb lower,
4370 void *cb_data);
4371 bool nir_shader_lower_instructions(nir_shader *shader,
4372 nir_instr_filter_cb filter,
4373 nir_lower_instr_cb lower,
4374 void *cb_data);
4375
4376 void nir_calc_dominance_impl(nir_function_impl *impl);
4377 void nir_calc_dominance(nir_shader *shader);
4378
4379 nir_block *nir_dominance_lca(nir_block *b1, nir_block *b2);
4380 bool nir_block_dominates(nir_block *parent, nir_block *child);
4381 bool nir_block_is_unreachable(nir_block *block);
4382
4383 void nir_dump_dom_tree_impl(nir_function_impl *impl, FILE *fp);
4384 void nir_dump_dom_tree(nir_shader *shader, FILE *fp);
4385
4386 void nir_dump_dom_frontier_impl(nir_function_impl *impl, FILE *fp);
4387 void nir_dump_dom_frontier(nir_shader *shader, FILE *fp);
4388
4389 void nir_dump_cfg_impl(nir_function_impl *impl, FILE *fp);
4390 void nir_dump_cfg(nir_shader *shader, FILE *fp);
4391
4392 void nir_gs_count_vertices_and_primitives(const nir_shader *shader,
4393 int *out_vtxcnt,
4394 int *out_prmcnt,
4395 unsigned num_streams);
4396
4397 typedef enum {
4398 nir_group_all,
4399 nir_group_same_resource_only,
4400 } nir_load_grouping;
4401
4402 void nir_group_loads(nir_shader *shader, nir_load_grouping grouping,
4403 unsigned max_distance);
4404
4405 bool nir_shrink_vec_array_vars(nir_shader *shader, nir_variable_mode modes);
4406 bool nir_split_array_vars(nir_shader *shader, nir_variable_mode modes);
4407 bool nir_split_var_copies(nir_shader *shader);
4408 bool nir_split_per_member_structs(nir_shader *shader);
4409 bool nir_split_struct_vars(nir_shader *shader, nir_variable_mode modes);
4410
4411 bool nir_lower_returns_impl(nir_function_impl *impl);
4412 bool nir_lower_returns(nir_shader *shader);
4413
4414 void nir_inline_function_impl(struct nir_builder *b,
4415 const nir_function_impl *impl,
4416 nir_ssa_def **params,
4417 struct hash_table *shader_var_remap);
4418 bool nir_inline_functions(nir_shader *shader);
4419
4420 void nir_find_inlinable_uniforms(nir_shader *shader);
4421 void nir_inline_uniforms(nir_shader *shader, unsigned num_uniforms,
4422 const uint32_t *uniform_values,
4423 const uint16_t *uniform_dw_offsets);
4424
4425 bool nir_propagate_invariant(nir_shader *shader, bool invariant_prim);
4426
4427 void nir_lower_var_copy_instr(nir_intrinsic_instr *copy, nir_shader *shader);
4428 void nir_lower_deref_copy_instr(struct nir_builder *b,
4429 nir_intrinsic_instr *copy);
4430 bool nir_lower_var_copies(nir_shader *shader);
4431
4432 bool nir_opt_memcpy(nir_shader *shader);
4433 bool nir_lower_memcpy(nir_shader *shader);
4434
4435 void nir_fixup_deref_modes(nir_shader *shader);
4436
4437 bool nir_lower_global_vars_to_local(nir_shader *shader);
4438
4439 typedef enum {
4440 nir_lower_direct_array_deref_of_vec_load = (1 << 0),
4441 nir_lower_indirect_array_deref_of_vec_load = (1 << 1),
4442 nir_lower_direct_array_deref_of_vec_store = (1 << 2),
4443 nir_lower_indirect_array_deref_of_vec_store = (1 << 3),
4444 } nir_lower_array_deref_of_vec_options;
4445
4446 bool nir_lower_array_deref_of_vec(nir_shader *shader, nir_variable_mode modes,
4447 nir_lower_array_deref_of_vec_options options);
4448
4449 bool nir_lower_indirect_derefs(nir_shader *shader, nir_variable_mode modes,
4450 uint32_t max_lower_array_len);
4451
4452 bool nir_lower_indirect_builtin_uniform_derefs(nir_shader *shader);
4453
4454 bool nir_lower_locals_to_regs(nir_shader *shader);
4455
4456 void nir_lower_io_to_temporaries(nir_shader *shader,
4457 nir_function_impl *entrypoint,
4458 bool outputs, bool inputs);
4459
4460 bool nir_lower_vars_to_scratch(nir_shader *shader,
4461 nir_variable_mode modes,
4462 int size_threshold,
4463 glsl_type_size_align_func size_align);
4464
4465 void nir_lower_clip_halfz(nir_shader *shader);
4466
4467 void nir_shader_gather_info(nir_shader *shader, nir_function_impl *entrypoint);
4468
4469 void nir_gather_ssa_types(nir_function_impl *impl,
4470 BITSET_WORD *float_types,
4471 BITSET_WORD *int_types);
4472
4473 void nir_assign_var_locations(nir_shader *shader, nir_variable_mode mode,
4474 unsigned *size,
4475 int (*type_size)(const struct glsl_type *, bool));
4476
4477 /* Some helpers to do very simple linking */
4478 bool nir_remove_unused_varyings(nir_shader *producer, nir_shader *consumer);
4479 bool nir_remove_unused_io_vars(nir_shader *shader, nir_variable_mode mode,
4480 uint64_t *used_by_other_stage,
4481 uint64_t *used_by_other_stage_patches);
4482 void nir_compact_varyings(nir_shader *producer, nir_shader *consumer,
4483 bool default_to_smooth_interp);
4484 void nir_link_xfb_varyings(nir_shader *producer, nir_shader *consumer);
4485 bool nir_link_opt_varyings(nir_shader *producer, nir_shader *consumer);
4486 void nir_link_varying_precision(nir_shader *producer, nir_shader *consumer);
4487
4488 bool nir_slot_is_sysval_output(gl_varying_slot slot);
4489 bool nir_slot_is_varying(gl_varying_slot slot);
4490 bool nir_slot_is_sysval_output_and_varying(gl_varying_slot slot);
4491 void nir_remove_varying(nir_intrinsic_instr *intr);
4492 void nir_remove_sysval_output(nir_intrinsic_instr *intr);
4493
4494 bool nir_lower_amul(nir_shader *shader,
4495 int (*type_size)(const struct glsl_type *, bool));
4496
4497 bool nir_lower_ubo_vec4(nir_shader *shader);
4498
4499 void nir_assign_io_var_locations(nir_shader *shader,
4500 nir_variable_mode mode,
4501 unsigned *size,
4502 gl_shader_stage stage);
4503
4504 typedef struct {
4505 uint8_t num_linked_io_vars;
4506 uint8_t num_linked_patch_io_vars;
4507 } nir_linked_io_var_info;
4508
4509 nir_linked_io_var_info
4510 nir_assign_linked_io_var_locations(nir_shader *producer,
4511 nir_shader *consumer);
4512
4513 typedef enum {
4514 /* If set, this causes all 64-bit IO operations to be lowered on-the-fly
4515 * to 32-bit operations. This is only valid for nir_var_shader_in/out
4516 * modes.
4517 */
4518 nir_lower_io_lower_64bit_to_32 = (1 << 0),
4519
4520 /* If set, this forces all non-flat fragment shader inputs to be
4521 * interpolated as if with the "sample" qualifier. This requires
4522 * nir_shader_compiler_options::use_interpolated_input_intrinsics.
4523 */
4524 nir_lower_io_force_sample_interpolation = (1 << 1),
4525 } nir_lower_io_options;
4526 bool nir_lower_io(nir_shader *shader,
4527 nir_variable_mode modes,
4528 int (*type_size)(const struct glsl_type *, bool),
4529 nir_lower_io_options);
4530
4531 bool nir_io_add_const_offset_to_base(nir_shader *nir, nir_variable_mode modes);
4532
4533 void
4534 nir_lower_io_passes(nir_shader *nir, struct nir_xfb_info *xfb);
4535
4536 bool
4537 nir_lower_vars_to_explicit_types(nir_shader *shader,
4538 nir_variable_mode modes,
4539 glsl_type_size_align_func type_info);
4540 void
4541 nir_gather_explicit_io_initializers(nir_shader *shader,
4542 void *dst, size_t dst_size,
4543 nir_variable_mode mode);
4544
4545 bool nir_lower_vec3_to_vec4(nir_shader *shader, nir_variable_mode modes);
4546
4547 typedef enum {
4548 /**
4549 * An address format which is a simple 32-bit global GPU address.
4550 */
4551 nir_address_format_32bit_global,
4552
4553 /**
4554 * An address format which is a simple 64-bit global GPU address.
4555 */
4556 nir_address_format_64bit_global,
4557
4558 /**
4559 * An address format which is a 64-bit global base address and a 32-bit
4560 * offset.
4561 *
4562 * The address is comprised as a 32-bit vec4 where .xy are a uint64_t base
4563 * address stored with the low bits in .x and high bits in .y, .z is
4564 * undefined, and .w is an offset. This is intended to match
4565 * 64bit_bounded_global but without the bounds checking.
4566 */
4567 nir_address_format_64bit_global_32bit_offset,
4568
4569 /**
4570 * An address format which is a bounds-checked 64-bit global GPU address.
4571 *
4572 * The address is comprised as a 32-bit vec4 where .xy are a uint64_t base
4573 * address stored with the low bits in .x and high bits in .y, .z is a
4574 * size, and .w is an offset. When the final I/O operation is lowered, .w
4575 * is checked against .z and the operation is predicated on the result.
4576 */
4577 nir_address_format_64bit_bounded_global,
4578
4579 /**
4580 * An address format which is comprised of a vec2 where the first
4581 * component is a buffer index and the second is an offset.
4582 */
4583 nir_address_format_32bit_index_offset,
4584
4585 /**
4586 * An address format which is a 64-bit value, where the high 32 bits
4587 * are a buffer index, and the low 32 bits are an offset.
4588 */
4589 nir_address_format_32bit_index_offset_pack64,
4590
4591 /**
4592 * An address format which is comprised of a vec3 where the first two
4593 * components specify the buffer and the third is an offset.
4594 */
4595 nir_address_format_vec2_index_32bit_offset,
4596
4597 /**
4598 * An address format which represents generic pointers with a 62-bit
4599 * pointer and a 2-bit enum in the top two bits. The top two bits have
4600 * the following meanings:
4601 *
4602 * - 0x0: Global memory
4603 * - 0x1: Shared memory
4604 * - 0x2: Scratch memory
4605 * - 0x3: Global memory
4606 *
4607 * The redundancy between 0x0 and 0x3 is because of Intel sign-extension of
4608 * addresses. Valid global memory addresses may naturally have either 0 or
4609 * ~0 as their high bits.
4610 *
4611 * Shared and scratch pointers are represented as 32-bit offsets with the
4612 * top 32 bits only being used for the enum. This allows us to avoid
4613 * 64-bit address calculations in a bunch of cases.
4614 */
4615 nir_address_format_62bit_generic,
4616
4617 /**
4618 * An address format which is a simple 32-bit offset.
4619 */
4620 nir_address_format_32bit_offset,
4621
4622 /**
4623 * An address format which is a simple 32-bit offset cast to 64-bit.
4624 */
4625 nir_address_format_32bit_offset_as_64bit,
4626
4627 /**
4628 * An address format representing a purely logical addressing model. In
4629 * this model, all deref chains must be complete from the dereference
4630 * operation to the variable. Cast derefs are not allowed. These
4631 * addresses will be 32-bit scalars but the format is immaterial because
4632 * you can always chase the chain.
4633 */
4634 nir_address_format_logical,
4635 } nir_address_format;
4636
4637 unsigned
4638 nir_address_format_bit_size(nir_address_format addr_format);
4639
4640 unsigned
4641 nir_address_format_num_components(nir_address_format addr_format);
4642
4643 static inline const struct glsl_type *
nir_address_format_to_glsl_type(nir_address_format addr_format)4644 nir_address_format_to_glsl_type(nir_address_format addr_format)
4645 {
4646 unsigned bit_size = nir_address_format_bit_size(addr_format);
4647 assert(bit_size == 32 || bit_size == 64);
4648 return glsl_vector_type(bit_size == 32 ? GLSL_TYPE_UINT : GLSL_TYPE_UINT64,
4649 nir_address_format_num_components(addr_format));
4650 }
4651
4652 const nir_const_value *nir_address_format_null_value(nir_address_format addr_format);
4653
4654 nir_ssa_def *nir_build_addr_ieq(struct nir_builder *b, nir_ssa_def *addr0, nir_ssa_def *addr1,
4655 nir_address_format addr_format);
4656
4657 nir_ssa_def *nir_build_addr_isub(struct nir_builder *b, nir_ssa_def *addr0, nir_ssa_def *addr1,
4658 nir_address_format addr_format);
4659
4660 nir_ssa_def * nir_explicit_io_address_from_deref(struct nir_builder *b,
4661 nir_deref_instr *deref,
4662 nir_ssa_def *base_addr,
4663 nir_address_format addr_format);
4664
4665 bool nir_get_explicit_deref_align(nir_deref_instr *deref,
4666 bool default_to_type_align,
4667 uint32_t *align_mul,
4668 uint32_t *align_offset);
4669
4670 void nir_lower_explicit_io_instr(struct nir_builder *b,
4671 nir_intrinsic_instr *io_instr,
4672 nir_ssa_def *addr,
4673 nir_address_format addr_format);
4674
4675 bool nir_lower_explicit_io(nir_shader *shader,
4676 nir_variable_mode modes,
4677 nir_address_format);
4678
4679 bool
4680 nir_lower_shader_calls(nir_shader *shader,
4681 nir_address_format address_format,
4682 unsigned stack_alignment,
4683 nir_shader ***resume_shaders_out,
4684 uint32_t *num_resume_shaders_out,
4685 void *mem_ctx);
4686
4687 nir_src *nir_get_io_offset_src(nir_intrinsic_instr *instr);
4688 nir_src *nir_get_io_arrayed_index_src(nir_intrinsic_instr *instr);
4689 nir_src *nir_get_shader_call_payload_src(nir_intrinsic_instr *call);
4690
4691 bool nir_is_arrayed_io(const nir_variable *var, gl_shader_stage stage);
4692
4693 bool nir_lower_regs_to_ssa_impl(nir_function_impl *impl);
4694 bool nir_lower_regs_to_ssa(nir_shader *shader);
4695 bool nir_lower_vars_to_ssa(nir_shader *shader);
4696
4697 bool nir_remove_dead_derefs(nir_shader *shader);
4698 bool nir_remove_dead_derefs_impl(nir_function_impl *impl);
4699
4700 typedef struct nir_remove_dead_variables_options {
4701 bool (*can_remove_var)(nir_variable *var, void *data);
4702 void *can_remove_var_data;
4703 } nir_remove_dead_variables_options;
4704
4705 bool nir_remove_dead_variables(nir_shader *shader, nir_variable_mode modes,
4706 const nir_remove_dead_variables_options *options);
4707
4708 bool nir_lower_variable_initializers(nir_shader *shader,
4709 nir_variable_mode modes);
4710 bool nir_zero_initialize_shared_memory(nir_shader *shader,
4711 const unsigned shared_size,
4712 const unsigned chunk_size);
4713
4714 bool nir_move_vec_src_uses_to_dest(nir_shader *shader);
4715 bool nir_lower_vec_to_movs(nir_shader *shader, nir_instr_writemask_filter_cb cb,
4716 const void *_data);
4717 void nir_lower_alpha_test(nir_shader *shader, enum compare_func func,
4718 bool alpha_to_one,
4719 const gl_state_index16 *alpha_ref_state_tokens);
4720 bool nir_lower_alu(nir_shader *shader);
4721
4722 bool nir_lower_flrp(nir_shader *shader, unsigned lowering_mask,
4723 bool always_precise);
4724
4725 bool nir_lower_alu_to_scalar(nir_shader *shader, nir_instr_filter_cb cb, const void *data);
4726 bool nir_lower_bool_to_bitsize(nir_shader *shader);
4727 bool nir_lower_bool_to_float(nir_shader *shader);
4728 bool nir_lower_bool_to_int32(nir_shader *shader);
4729 bool nir_opt_simplify_convert_alu_types(nir_shader *shader);
4730 bool nir_lower_convert_alu_types(nir_shader *shader,
4731 bool (*should_lower)(nir_intrinsic_instr *));
4732 bool nir_lower_constant_convert_alu_types(nir_shader *shader);
4733 bool nir_lower_alu_conversion_to_intrinsic(nir_shader *shader);
4734 bool nir_lower_int_to_float(nir_shader *shader);
4735 bool nir_lower_load_const_to_scalar(nir_shader *shader);
4736 bool nir_lower_read_invocation_to_scalar(nir_shader *shader);
4737 bool nir_lower_phis_to_scalar(nir_shader *shader, bool lower_all);
4738 void nir_lower_io_arrays_to_elements(nir_shader *producer, nir_shader *consumer);
4739 void nir_lower_io_arrays_to_elements_no_indirects(nir_shader *shader,
4740 bool outputs_only);
4741 void nir_lower_io_to_scalar(nir_shader *shader, nir_variable_mode mask);
4742 bool nir_lower_io_to_scalar_early(nir_shader *shader, nir_variable_mode mask);
4743 bool nir_lower_io_to_vector(nir_shader *shader, nir_variable_mode mask);
4744 bool nir_vectorize_tess_levels(nir_shader *shader);
4745
4746 bool nir_lower_fragcolor(nir_shader *shader, unsigned max_cbufs);
4747 bool nir_lower_fragcoord_wtrans(nir_shader *shader);
4748 void nir_lower_viewport_transform(nir_shader *shader);
4749 bool nir_lower_uniforms_to_ubo(nir_shader *shader, bool dword_packed, bool load_vec4);
4750
4751 bool nir_lower_is_helper_invocation(nir_shader *shader);
4752
4753 typedef struct nir_lower_subgroups_options {
4754 uint8_t subgroup_size;
4755 uint8_t ballot_bit_size;
4756 uint8_t ballot_components;
4757 bool lower_to_scalar:1;
4758 bool lower_vote_trivial:1;
4759 bool lower_vote_eq:1;
4760 bool lower_subgroup_masks:1;
4761 bool lower_relative_shuffle:1;
4762 bool lower_shuffle_to_32bit:1;
4763 bool lower_shuffle_to_swizzle_amd:1;
4764 bool lower_shuffle:1;
4765 bool lower_quad:1;
4766 bool lower_quad_broadcast_dynamic:1;
4767 bool lower_quad_broadcast_dynamic_to_const:1;
4768 bool lower_elect:1;
4769 bool lower_read_invocation_to_cond:1;
4770 } nir_lower_subgroups_options;
4771
4772 bool nir_lower_subgroups(nir_shader *shader,
4773 const nir_lower_subgroups_options *options);
4774
4775 bool nir_lower_system_values(nir_shader *shader);
4776
4777 typedef struct nir_lower_compute_system_values_options {
4778 bool has_base_global_invocation_id:1;
4779 bool has_base_workgroup_id:1;
4780 bool shuffle_local_ids_for_quad_derivatives:1;
4781 bool lower_local_invocation_index:1;
4782 bool lower_cs_local_id_to_index:1;
4783 bool lower_workgroup_id_to_index:1;
4784 } nir_lower_compute_system_values_options;
4785
4786 bool nir_lower_compute_system_values(nir_shader *shader,
4787 const nir_lower_compute_system_values_options *options);
4788
4789 struct nir_lower_sysvals_to_varyings_options {
4790 bool frag_coord:1;
4791 bool front_face:1;
4792 bool point_coord:1;
4793 };
4794
4795 bool
4796 nir_lower_sysvals_to_varyings(nir_shader *shader,
4797 const struct nir_lower_sysvals_to_varyings_options *options);
4798
4799 enum PACKED nir_lower_tex_packing {
4800 /** No packing */
4801 nir_lower_tex_packing_none = 0,
4802 /**
4803 * The sampler returns up to 2 32-bit words of half floats or 16-bit signed
4804 * or unsigned ints based on the sampler type
4805 */
4806 nir_lower_tex_packing_16,
4807 /** The sampler returns 1 32-bit word of 4x8 unorm */
4808 nir_lower_tex_packing_8,
4809 };
4810
4811 typedef struct nir_lower_tex_options {
4812 /**
4813 * bitmask of (1 << GLSL_SAMPLER_DIM_x) to control for which
4814 * sampler types a texture projector is lowered.
4815 */
4816 unsigned lower_txp;
4817
4818 /**
4819 * If true, lower texture projector for any array sampler dims
4820 */
4821 bool lower_txp_array;
4822
4823 /**
4824 * If true, lower away nir_tex_src_offset for all texelfetch instructions.
4825 */
4826 bool lower_txf_offset;
4827
4828 /**
4829 * If true, lower away nir_tex_src_offset for all rect textures.
4830 */
4831 bool lower_rect_offset;
4832
4833 /**
4834 * If not NULL, this filter will return true for tex instructions that
4835 * should lower away nir_tex_src_offset.
4836 */
4837 nir_instr_filter_cb lower_offset_filter;
4838
4839 /**
4840 * If true, lower rect textures to 2D, using txs to fetch the
4841 * texture dimensions and dividing the texture coords by the
4842 * texture dims to normalize.
4843 */
4844 bool lower_rect;
4845
4846 /**
4847 * If true, convert yuv to rgb.
4848 */
4849 unsigned lower_y_uv_external;
4850 unsigned lower_y_u_v_external;
4851 unsigned lower_yx_xuxv_external;
4852 unsigned lower_xy_uxvx_external;
4853 unsigned lower_ayuv_external;
4854 unsigned lower_xyuv_external;
4855 unsigned lower_yuv_external;
4856 unsigned lower_yu_yv_external;
4857 unsigned lower_y41x_external;
4858 unsigned bt709_external;
4859 unsigned bt2020_external;
4860
4861 /**
4862 * To emulate certain texture wrap modes, this can be used
4863 * to saturate the specified tex coord to [0.0, 1.0]. The
4864 * bits are according to sampler #, ie. if, for example:
4865 *
4866 * (conf->saturate_s & (1 << n))
4867 *
4868 * is true, then the s coord for sampler n is saturated.
4869 *
4870 * Note that clamping must happen *after* projector lowering
4871 * so any projected texture sample instruction with a clamped
4872 * coordinate gets automatically lowered, regardless of the
4873 * 'lower_txp' setting.
4874 */
4875 unsigned saturate_s;
4876 unsigned saturate_t;
4877 unsigned saturate_r;
4878
4879 /* Bitmask of textures that need swizzling.
4880 *
4881 * If (swizzle_result & (1 << texture_index)), then the swizzle in
4882 * swizzles[texture_index] is applied to the result of the texturing
4883 * operation.
4884 */
4885 unsigned swizzle_result;
4886
4887 /* A swizzle for each texture. Values 0-3 represent x, y, z, or w swizzles
4888 * while 4 and 5 represent 0 and 1 respectively.
4889 *
4890 * Indexed by texture-id.
4891 */
4892 uint8_t swizzles[32][4];
4893
4894 /* Can be used to scale sampled values in range required by the
4895 * format.
4896 *
4897 * Indexed by texture-id.
4898 */
4899 float scale_factors[32];
4900
4901 /**
4902 * Bitmap of textures that need srgb to linear conversion. If
4903 * (lower_srgb & (1 << texture_index)) then the rgb (xyz) components
4904 * of the texture are lowered to linear.
4905 */
4906 unsigned lower_srgb;
4907
4908 /**
4909 * If true, lower nir_texop_txd on cube maps with nir_texop_txl.
4910 */
4911 bool lower_txd_cube_map;
4912
4913 /**
4914 * If true, lower nir_texop_txd on 3D surfaces with nir_texop_txl.
4915 */
4916 bool lower_txd_3d;
4917
4918 /**
4919 * If true, lower nir_texop_txd on shadow samplers (except cube maps)
4920 * with nir_texop_txl. Notice that cube map shadow samplers are lowered
4921 * with lower_txd_cube_map.
4922 */
4923 bool lower_txd_shadow;
4924
4925 /**
4926 * If true, lower nir_texop_txd on all samplers to a nir_texop_txl.
4927 * Implies lower_txd_cube_map and lower_txd_shadow.
4928 */
4929 bool lower_txd;
4930
4931 /**
4932 * If true, lower nir_texop_txb that try to use shadow compare and min_lod
4933 * at the same time to a nir_texop_lod, some math, and nir_texop_tex.
4934 */
4935 bool lower_txb_shadow_clamp;
4936
4937 /**
4938 * If true, lower nir_texop_txd on shadow samplers when it uses min_lod
4939 * with nir_texop_txl. This includes cube maps.
4940 */
4941 bool lower_txd_shadow_clamp;
4942
4943 /**
4944 * If true, lower nir_texop_txd on when it uses both offset and min_lod
4945 * with nir_texop_txl. This includes cube maps.
4946 */
4947 bool lower_txd_offset_clamp;
4948
4949 /**
4950 * If true, lower nir_texop_txd with min_lod to a nir_texop_txl if the
4951 * sampler is bindless.
4952 */
4953 bool lower_txd_clamp_bindless_sampler;
4954
4955 /**
4956 * If true, lower nir_texop_txd with min_lod to a nir_texop_txl if the
4957 * sampler index is not statically determinable to be less than 16.
4958 */
4959 bool lower_txd_clamp_if_sampler_index_not_lt_16;
4960
4961 /**
4962 * If true, lower nir_texop_txs with a non-0-lod into nir_texop_txs with
4963 * 0-lod followed by a nir_ishr.
4964 */
4965 bool lower_txs_lod;
4966
4967 /**
4968 * If true, lower nir_texop_txs for cube arrays to a nir_texop_txs with a
4969 * 2D array type followed by a nir_idiv by 6.
4970 */
4971 bool lower_txs_cube_array;
4972
4973 /**
4974 * If true, apply a .bagr swizzle on tg4 results to handle Broadcom's
4975 * mixed-up tg4 locations.
4976 */
4977 bool lower_tg4_broadcom_swizzle;
4978
4979 /**
4980 * If true, lowers tg4 with 4 constant offsets to 4 tg4 calls
4981 */
4982 bool lower_tg4_offsets;
4983
4984 /**
4985 * Lower txf_ms to fragment_mask_fetch and fragment_fetch and samples_identical to
4986 * fragment_mask_fetch.
4987 */
4988 bool lower_to_fragment_fetch_amd;
4989
4990 /**
4991 * To lower packed sampler return formats.
4992 *
4993 * Indexed by sampler-id.
4994 */
4995 enum nir_lower_tex_packing lower_tex_packing[32];
4996
4997 /**
4998 * If true, lower nir_texop_lod to return -FLT_MAX if the sum of the
4999 * absolute values of derivatives is 0 for all coordinates.
5000 */
5001 bool lower_lod_zero_width;
5002
5003 /**
5004 * Payload data to be sent to callback / filter functions.
5005 */
5006 void *callback_data;
5007 } nir_lower_tex_options;
5008
5009 /** Lowers complex texture instructions to simpler ones */
5010 bool nir_lower_tex(nir_shader *shader,
5011 const nir_lower_tex_options *options);
5012
5013
5014 typedef struct nir_lower_tex_shadow_swizzle {
5015 unsigned swizzle_r:3;
5016 unsigned swizzle_g:3;
5017 unsigned swizzle_b:3;
5018 unsigned swizzle_a:3;
5019 } nir_lower_tex_shadow_swizzle;
5020
5021 bool
5022 nir_lower_tex_shadow(nir_shader *s,
5023 unsigned n_states,
5024 enum compare_func *compare_func,
5025 nir_lower_tex_shadow_swizzle *tex_swizzles);
5026
5027 typedef struct nir_lower_image_options {
5028 /**
5029 * If true, lower cube size operations.
5030 */
5031 bool lower_cube_size;
5032 } nir_lower_image_options;
5033
5034 bool nir_lower_image(nir_shader *nir,
5035 const nir_lower_image_options *options);
5036
5037 bool nir_lower_readonly_images_to_tex(nir_shader *shader, bool per_variable);
5038
5039 enum nir_lower_non_uniform_access_type {
5040 nir_lower_non_uniform_ubo_access = (1 << 0),
5041 nir_lower_non_uniform_ssbo_access = (1 << 1),
5042 nir_lower_non_uniform_texture_access = (1 << 2),
5043 nir_lower_non_uniform_image_access = (1 << 3),
5044 };
5045
5046 /* Given the nir_src used for the resource, return the channels which might be non-uniform. */
5047 typedef nir_component_mask_t (*nir_lower_non_uniform_access_callback)(const nir_src *, void *);
5048
5049 typedef struct nir_lower_non_uniform_access_options {
5050 enum nir_lower_non_uniform_access_type types;
5051 nir_lower_non_uniform_access_callback callback;
5052 void *callback_data;
5053 } nir_lower_non_uniform_access_options;
5054
5055 bool nir_lower_non_uniform_access(nir_shader *shader,
5056 const nir_lower_non_uniform_access_options *options);
5057
5058 typedef struct {
5059 /* If true, a 32-bit division lowering based on NV50LegalizeSSA::handleDIV()
5060 * is used. It is the faster of the two but it is not exact in some cases
5061 * (for example, 1091317713u / 1034u gives 5209173 instead of 1055432).
5062 *
5063 * If false, a lowering based on AMDGPUTargetLowering::LowerUDIVREM() and
5064 * AMDGPUTargetLowering::LowerSDIVREM() is used. It requires more
5065 * instructions than the nv50 path and many of them are integer
5066 * multiplications, so it is probably slower. It should always return the
5067 * correct result, though.
5068 */
5069 bool imprecise_32bit_lowering;
5070
5071 /* Whether 16-bit floating point arithmetic should be allowed in 8-bit
5072 * division lowering
5073 */
5074 bool allow_fp16;
5075 } nir_lower_idiv_options;
5076
5077 bool nir_lower_idiv(nir_shader *shader, const nir_lower_idiv_options *options);
5078
5079 typedef struct nir_input_attachment_options {
5080 bool use_fragcoord_sysval;
5081 bool use_layer_id_sysval;
5082 bool use_view_id_for_layer;
5083 } nir_input_attachment_options;
5084
5085 bool nir_lower_input_attachments(nir_shader *shader,
5086 const nir_input_attachment_options *options);
5087
5088 bool nir_lower_clip_vs(nir_shader *shader, unsigned ucp_enables,
5089 bool use_vars,
5090 bool use_clipdist_array,
5091 const gl_state_index16 clipplane_state_tokens[][STATE_LENGTH]);
5092 bool nir_lower_clip_gs(nir_shader *shader, unsigned ucp_enables,
5093 bool use_clipdist_array,
5094 const gl_state_index16 clipplane_state_tokens[][STATE_LENGTH]);
5095 bool nir_lower_clip_fs(nir_shader *shader, unsigned ucp_enables,
5096 bool use_clipdist_array);
5097 bool nir_lower_clip_cull_distance_arrays(nir_shader *nir);
5098 bool nir_lower_clip_disable(nir_shader *shader, unsigned clip_plane_enable);
5099
5100 void nir_lower_point_size_mov(nir_shader *shader,
5101 const gl_state_index16 *pointsize_state_tokens);
5102
5103 bool nir_lower_frexp(nir_shader *nir);
5104
5105 void nir_lower_two_sided_color(nir_shader *shader, bool face_sysval);
5106
5107 bool nir_lower_clamp_color_outputs(nir_shader *shader);
5108
5109 bool nir_lower_flatshade(nir_shader *shader);
5110
5111 void nir_lower_passthrough_edgeflags(nir_shader *shader);
5112 bool nir_lower_patch_vertices(nir_shader *nir, unsigned static_count,
5113 const gl_state_index16 *uniform_state_tokens);
5114
5115 typedef struct nir_lower_wpos_ytransform_options {
5116 gl_state_index16 state_tokens[STATE_LENGTH];
5117 bool fs_coord_origin_upper_left :1;
5118 bool fs_coord_origin_lower_left :1;
5119 bool fs_coord_pixel_center_integer :1;
5120 bool fs_coord_pixel_center_half_integer :1;
5121 } nir_lower_wpos_ytransform_options;
5122
5123 bool nir_lower_wpos_ytransform(nir_shader *shader,
5124 const nir_lower_wpos_ytransform_options *options);
5125 bool nir_lower_wpos_center(nir_shader *shader);
5126
5127 bool nir_lower_pntc_ytransform(nir_shader *shader,
5128 const gl_state_index16 clipplane_state_tokens[][STATE_LENGTH]);
5129
5130 bool nir_lower_wrmasks(nir_shader *shader, nir_instr_filter_cb cb, const void *data);
5131
5132 bool nir_lower_fb_read(nir_shader *shader);
5133
5134 typedef struct nir_lower_drawpixels_options {
5135 gl_state_index16 texcoord_state_tokens[STATE_LENGTH];
5136 gl_state_index16 scale_state_tokens[STATE_LENGTH];
5137 gl_state_index16 bias_state_tokens[STATE_LENGTH];
5138 unsigned drawpix_sampler;
5139 unsigned pixelmap_sampler;
5140 bool pixel_maps :1;
5141 bool scale_and_bias :1;
5142 } nir_lower_drawpixels_options;
5143
5144 void nir_lower_drawpixels(nir_shader *shader,
5145 const nir_lower_drawpixels_options *options);
5146
5147 typedef struct nir_lower_bitmap_options {
5148 unsigned sampler;
5149 bool swizzle_xxxx;
5150 } nir_lower_bitmap_options;
5151
5152 void nir_lower_bitmap(nir_shader *shader, const nir_lower_bitmap_options *options);
5153
5154 bool nir_lower_atomics_to_ssbo(nir_shader *shader);
5155
5156 typedef enum {
5157 nir_lower_int_source_mods = 1 << 0,
5158 nir_lower_fabs_source_mods = 1 << 1,
5159 nir_lower_fneg_source_mods = 1 << 2,
5160 nir_lower_64bit_source_mods = 1 << 3,
5161 nir_lower_triop_abs = 1 << 4,
5162 nir_lower_all_source_mods = (1 << 5) - 1
5163 } nir_lower_to_source_mods_flags;
5164
5165 #define nir_lower_float_source_mods (nir_lower_fabs_source_mods | nir_lower_fneg_source_mods)
5166
5167 bool nir_lower_to_source_mods(nir_shader *shader, nir_lower_to_source_mods_flags options);
5168
5169 typedef enum {
5170 nir_lower_gs_intrinsics_per_stream = 1 << 0,
5171 nir_lower_gs_intrinsics_count_primitives = 1 << 1,
5172 nir_lower_gs_intrinsics_count_vertices_per_primitive = 1 << 2,
5173 nir_lower_gs_intrinsics_overwrite_incomplete = 1 << 3,
5174 } nir_lower_gs_intrinsics_flags;
5175
5176 bool nir_lower_gs_intrinsics(nir_shader *shader, nir_lower_gs_intrinsics_flags options);
5177
5178 typedef unsigned (*nir_lower_bit_size_callback)(const nir_instr *, void *);
5179
5180 bool nir_lower_bit_size(nir_shader *shader,
5181 nir_lower_bit_size_callback callback,
5182 void *callback_data);
5183 bool nir_lower_64bit_phis(nir_shader *shader);
5184
5185 nir_lower_int64_options nir_lower_int64_op_to_options_mask(nir_op opcode);
5186 bool nir_lower_int64(nir_shader *shader);
5187
5188 nir_lower_doubles_options nir_lower_doubles_op_to_options_mask(nir_op opcode);
5189 bool nir_lower_doubles(nir_shader *shader, const nir_shader *softfp64,
5190 nir_lower_doubles_options options);
5191 bool nir_lower_pack(nir_shader *shader);
5192
5193 bool nir_recompute_io_bases(nir_shader *nir, nir_variable_mode modes);
5194 bool nir_lower_mediump_io(nir_shader *nir, nir_variable_mode modes,
5195 uint64_t varying_mask, bool use_16bit_slots);
5196 bool nir_force_mediump_io(nir_shader *nir, nir_variable_mode modes,
5197 nir_alu_type types);
5198 bool nir_unpack_16bit_varying_slots(nir_shader *nir, nir_variable_mode modes);
5199 bool nir_fold_16bit_sampler_conversions(nir_shader *nir,
5200 unsigned tex_src_types);
5201
5202 typedef struct {
5203 bool legalize_type; /* whether this src should be legalized */
5204 uint8_t bit_size; /* bit_size to enforce */
5205 nir_tex_src_type match_src; /* if bit_size is 0, match bit size of this */
5206 } nir_tex_src_type_constraint, nir_tex_src_type_constraints[nir_num_tex_src_types];
5207
5208 bool nir_legalize_16bit_sampler_srcs(nir_shader *nir,
5209 nir_tex_src_type_constraints constraints);
5210
5211 bool nir_lower_point_size(nir_shader *shader, float min, float max);
5212
5213 void nir_lower_texcoord_replace(nir_shader *s, unsigned coord_replace,
5214 bool point_coord_is_sysval, bool yinvert);
5215
5216 typedef enum {
5217 nir_lower_interpolation_at_sample = (1 << 1),
5218 nir_lower_interpolation_at_offset = (1 << 2),
5219 nir_lower_interpolation_centroid = (1 << 3),
5220 nir_lower_interpolation_pixel = (1 << 4),
5221 nir_lower_interpolation_sample = (1 << 5),
5222 } nir_lower_interpolation_options;
5223
5224 bool nir_lower_interpolation(nir_shader *shader,
5225 nir_lower_interpolation_options options);
5226
5227 bool nir_lower_discard_if(nir_shader *shader);
5228
5229 bool nir_lower_discard_or_demote(nir_shader *shader,
5230 bool force_correct_quad_ops_after_discard);
5231
5232 bool nir_lower_memory_model(nir_shader *shader);
5233
5234 bool nir_lower_goto_ifs(nir_shader *shader);
5235
5236 bool nir_shader_uses_view_index(nir_shader *shader);
5237 bool nir_can_lower_multiview(nir_shader *shader);
5238 bool nir_lower_multiview(nir_shader *shader, uint32_t view_mask);
5239
5240
5241 bool nir_lower_fp16_casts(nir_shader *shader);
5242 bool nir_normalize_cubemap_coords(nir_shader *shader);
5243
5244 bool nir_shader_supports_implicit_lod(nir_shader *shader);
5245
5246 void nir_live_ssa_defs_impl(nir_function_impl *impl);
5247
5248 const BITSET_WORD *nir_get_live_ssa_defs(nir_cursor cursor, void *mem_ctx);
5249
5250 void nir_loop_analyze_impl(nir_function_impl *impl,
5251 nir_variable_mode indirect_mask);
5252
5253 bool nir_ssa_defs_interfere(nir_ssa_def *a, nir_ssa_def *b);
5254
5255 bool nir_repair_ssa_impl(nir_function_impl *impl);
5256 bool nir_repair_ssa(nir_shader *shader);
5257
5258 void nir_convert_loop_to_lcssa(nir_loop *loop);
5259 bool nir_convert_to_lcssa(nir_shader *shader, bool skip_invariants, bool skip_bool_invariants);
5260 void nir_divergence_analysis(nir_shader *shader);
5261 bool nir_update_instr_divergence(nir_shader *shader, nir_instr *instr);
5262 bool nir_has_divergent_loop(nir_shader *shader);
5263
5264 /* If phi_webs_only is true, only convert SSA values involved in phi nodes to
5265 * registers. If false, convert all values (even those not involved in a phi
5266 * node) to registers.
5267 */
5268 bool nir_convert_from_ssa(nir_shader *shader, bool phi_webs_only);
5269
5270 bool nir_lower_phis_to_regs_block(nir_block *block);
5271 bool nir_lower_ssa_defs_to_regs_block(nir_block *block);
5272 bool nir_rematerialize_derefs_in_use_blocks_impl(nir_function_impl *impl);
5273
5274 bool nir_lower_samplers(nir_shader *shader);
5275 bool nir_lower_ssbo(nir_shader *shader);
5276
5277 typedef struct nir_lower_printf_options {
5278 bool treat_doubles_as_floats : 1;
5279 unsigned max_buffer_size;
5280 } nir_lower_printf_options;
5281
5282 bool nir_lower_printf(nir_shader *nir, const nir_lower_printf_options *options);
5283
5284 /* This is here for unit tests. */
5285 bool nir_opt_comparison_pre_impl(nir_function_impl *impl);
5286
5287 bool nir_opt_comparison_pre(nir_shader *shader);
5288
5289 typedef struct nir_opt_access_options {
5290 bool is_vulkan;
5291 bool infer_non_readable;
5292 } nir_opt_access_options;
5293
5294 bool nir_opt_access(nir_shader *shader, const nir_opt_access_options *options);
5295 bool nir_opt_algebraic(nir_shader *shader);
5296 bool nir_opt_algebraic_before_ffma(nir_shader *shader);
5297 bool nir_opt_algebraic_late(nir_shader *shader);
5298 bool nir_opt_algebraic_distribute_src_mods(nir_shader *shader);
5299 bool nir_opt_constant_folding(nir_shader *shader);
5300
5301 /* Try to combine a and b into a. Return true if combination was possible,
5302 * which will result in b being removed by the pass. Return false if
5303 * combination wasn't possible.
5304 */
5305 typedef bool (*nir_combine_memory_barrier_cb)(
5306 nir_intrinsic_instr *a, nir_intrinsic_instr *b, void *data);
5307
5308 bool nir_opt_combine_memory_barriers(nir_shader *shader,
5309 nir_combine_memory_barrier_cb combine_cb,
5310 void *data);
5311
5312 bool nir_opt_combine_stores(nir_shader *shader, nir_variable_mode modes);
5313
5314 bool nir_copy_prop_impl(nir_function_impl *impl);
5315 bool nir_copy_prop(nir_shader *shader);
5316
5317 bool nir_opt_copy_prop_vars(nir_shader *shader);
5318
5319 bool nir_opt_cse(nir_shader *shader);
5320
5321 bool nir_opt_dce(nir_shader *shader);
5322
5323 bool nir_opt_dead_cf(nir_shader *shader);
5324
5325 bool nir_opt_dead_write_vars(nir_shader *shader);
5326
5327 bool nir_opt_deref_impl(nir_function_impl *impl);
5328 bool nir_opt_deref(nir_shader *shader);
5329
5330 bool nir_opt_find_array_copies(nir_shader *shader);
5331
5332 bool nir_opt_fragdepth(nir_shader *shader);
5333
5334 bool nir_opt_gcm(nir_shader *shader, bool value_number);
5335
5336 bool nir_opt_idiv_const(nir_shader *shader, unsigned min_bit_size);
5337
5338 bool nir_opt_if(nir_shader *shader, bool aggressive_last_continue);
5339
5340 bool nir_opt_intrinsics(nir_shader *shader);
5341
5342 bool nir_opt_large_constants(nir_shader *shader,
5343 glsl_type_size_align_func size_align,
5344 unsigned threshold);
5345
5346 bool nir_opt_loop_unroll(nir_shader *shader);
5347
5348 typedef enum {
5349 nir_move_const_undef = (1 << 0),
5350 nir_move_load_ubo = (1 << 1),
5351 nir_move_load_input = (1 << 2),
5352 nir_move_comparisons = (1 << 3),
5353 nir_move_copies = (1 << 4),
5354 nir_move_load_ssbo = (1 << 5),
5355 nir_move_load_uniform = (1 << 6),
5356 } nir_move_options;
5357
5358 bool nir_can_move_instr(nir_instr *instr, nir_move_options options);
5359
5360 bool nir_opt_sink(nir_shader *shader, nir_move_options options);
5361
5362 bool nir_opt_move(nir_shader *shader, nir_move_options options);
5363
5364 typedef struct {
5365 /** nir_load_uniform max base offset */
5366 uint32_t uniform_max;
5367
5368 /** nir_load_ubo_vec4 max base offset */
5369 uint32_t ubo_vec4_max;
5370
5371 /** nir_var_mem_shared max base offset */
5372 uint32_t shared_max;
5373
5374 /** nir_load/store_buffer_amd max base offset */
5375 uint32_t buffer_max;
5376 } nir_opt_offsets_options;
5377
5378 bool nir_opt_offsets(nir_shader *shader, const nir_opt_offsets_options *options);
5379
5380 bool nir_opt_peephole_select(nir_shader *shader, unsigned limit,
5381 bool indirect_load_ok, bool expensive_alu_ok);
5382
5383 bool nir_opt_rematerialize_compares(nir_shader *shader);
5384
5385 bool nir_opt_remove_phis(nir_shader *shader);
5386 bool nir_opt_remove_phis_block(nir_block *block);
5387
5388 bool nir_opt_phi_precision(nir_shader *shader);
5389
5390 bool nir_opt_shrink_stores(nir_shader *shader, bool shrink_image_store);
5391
5392 bool nir_opt_shrink_vectors(nir_shader *shader);
5393
5394 bool nir_opt_trivial_continues(nir_shader *shader);
5395
5396 bool nir_opt_undef(nir_shader *shader);
5397
5398 bool nir_lower_undef_to_zero(nir_shader *shader);
5399
5400 bool nir_opt_uniform_atomics(nir_shader *shader);
5401
5402 typedef bool (*nir_opt_vectorize_cb)(const nir_instr *instr, void *data);
5403
5404 bool nir_opt_vectorize(nir_shader *shader, nir_opt_vectorize_cb filter,
5405 void *data);
5406
5407 bool nir_opt_conditional_discard(nir_shader *shader);
5408 bool nir_opt_move_discards_to_top(nir_shader *shader);
5409
5410 bool nir_opt_ray_queries(nir_shader *shader);
5411
5412 typedef bool (*nir_should_vectorize_mem_func)(unsigned align_mul,
5413 unsigned align_offset,
5414 unsigned bit_size,
5415 unsigned num_components,
5416 nir_intrinsic_instr *low, nir_intrinsic_instr *high,
5417 void *data);
5418
5419 typedef struct {
5420 nir_should_vectorize_mem_func callback;
5421 nir_variable_mode modes;
5422 nir_variable_mode robust_modes;
5423 void *cb_data;
5424 } nir_load_store_vectorize_options;
5425
5426 bool nir_opt_load_store_vectorize(nir_shader *shader, const nir_load_store_vectorize_options *options);
5427
5428 void nir_sweep(nir_shader *shader);
5429
5430 void nir_remap_dual_slot_attributes(nir_shader *shader,
5431 uint64_t *dual_slot_inputs);
5432 uint64_t nir_get_single_slot_attribs_mask(uint64_t attribs, uint64_t dual_slot);
5433
5434 nir_intrinsic_op nir_intrinsic_from_system_value(gl_system_value val);
5435 gl_system_value nir_system_value_from_intrinsic(nir_intrinsic_op intrin);
5436
5437 static inline bool
nir_variable_is_in_ubo(const nir_variable * var)5438 nir_variable_is_in_ubo(const nir_variable *var)
5439 {
5440 return (var->data.mode == nir_var_mem_ubo &&
5441 var->interface_type != NULL);
5442 }
5443
5444 static inline bool
nir_variable_is_in_ssbo(const nir_variable * var)5445 nir_variable_is_in_ssbo(const nir_variable *var)
5446 {
5447 return (var->data.mode == nir_var_mem_ssbo &&
5448 var->interface_type != NULL);
5449 }
5450
5451 static inline bool
nir_variable_is_in_block(const nir_variable * var)5452 nir_variable_is_in_block(const nir_variable *var)
5453 {
5454 return nir_variable_is_in_ubo(var) || nir_variable_is_in_ssbo(var);
5455 }
5456
5457 typedef struct nir_unsigned_upper_bound_config {
5458 unsigned min_subgroup_size;
5459 unsigned max_subgroup_size;
5460 unsigned max_workgroup_invocations;
5461 unsigned max_workgroup_count[3];
5462 unsigned max_workgroup_size[3];
5463
5464 uint32_t vertex_attrib_max[32];
5465 } nir_unsigned_upper_bound_config;
5466
5467 uint32_t
5468 nir_unsigned_upper_bound(nir_shader *shader, struct hash_table *range_ht,
5469 nir_ssa_scalar scalar,
5470 const nir_unsigned_upper_bound_config *config);
5471
5472 bool
5473 nir_addition_might_overflow(nir_shader *shader, struct hash_table *range_ht,
5474 nir_ssa_scalar ssa, unsigned const_val,
5475 const nir_unsigned_upper_bound_config *config);
5476
5477 typedef enum {
5478 nir_ray_query_value_intersection_type,
5479 nir_ray_query_value_intersection_t,
5480 nir_ray_query_value_intersection_instance_custom_index,
5481 nir_ray_query_value_intersection_instance_id,
5482 nir_ray_query_value_intersection_instance_sbt_index,
5483 nir_ray_query_value_intersection_geometry_index,
5484 nir_ray_query_value_intersection_primitive_index,
5485 nir_ray_query_value_intersection_barycentrics,
5486 nir_ray_query_value_intersection_front_face,
5487 nir_ray_query_value_intersection_object_ray_direction,
5488 nir_ray_query_value_intersection_object_ray_origin,
5489 nir_ray_query_value_intersection_object_to_world,
5490 nir_ray_query_value_intersection_world_to_object,
5491 nir_ray_query_value_intersection_candidate_aabb_opaque,
5492 nir_ray_query_value_tmin,
5493 nir_ray_query_value_flags,
5494 nir_ray_query_value_world_ray_direction,
5495 nir_ray_query_value_world_ray_origin,
5496 } nir_ray_query_value;
5497
5498 typedef struct {
5499 /* True if gl_DrawID is considered uniform, i.e. if the preamble is run
5500 * at least once per "internal" draw rather than per user-visible draw.
5501 */
5502 bool drawid_uniform;
5503
5504 /* True if the subgroup size is uniform. */
5505 bool subgroup_size_uniform;
5506
5507 /* size/align for load/store_preamble. */
5508 void (*def_size)(nir_ssa_def *def, unsigned *size, unsigned *align);
5509
5510 /* Total available size for load/store_preamble storage, in units
5511 * determined by def_size.
5512 */
5513 unsigned preamble_storage_size;
5514
5515 /* Give the cost for an instruction. nir_opt_preamble will prioritize
5516 * instructions with higher costs. Instructions with cost 0 may still be
5517 * lifted, but only when required to lift other instructions with non-0
5518 * cost (e.g. a load_const source of an expression).
5519 */
5520 float (*instr_cost_cb)(nir_instr *instr, const void *data);
5521
5522 /* Give the cost of rewriting the instruction to use load_preamble. This
5523 * may happen from inserting move instructions, etc. If the benefit doesn't
5524 * exceed the cost here then we won't rewrite it.
5525 */
5526 float (*rewrite_cost_cb)(nir_ssa_def *def, const void *data);
5527
5528 /* Instructions whose definitions should not be rewritten. These could
5529 * still be moved to the preamble, but they shouldn't be the root of a
5530 * replacement expression. Instructions with cost 0 and derefs are
5531 * automatically included by the pass.
5532 */
5533 nir_instr_filter_cb avoid_instr_cb;
5534
5535 const void *cb_data;
5536 } nir_opt_preamble_options;
5537
5538 bool
5539 nir_opt_preamble(nir_shader *shader,
5540 const nir_opt_preamble_options *options,
5541 unsigned *size);
5542
5543 nir_function_impl *nir_shader_get_preamble(nir_shader *shader);
5544
5545 #include "nir_inline_helpers.h"
5546
5547 #ifdef __cplusplus
5548 } /* extern "C" */
5549 #endif
5550
5551 #endif /* NIR_H */
5552