1/* Copyright (C) 2018 Red Hat
2 *
3 * Permission is hereby granted, free of charge, to any person obtaining a
4 * copy of this software and associated documentation files (the "Software"),
5 * to deal in the Software without restriction, including without limitation
6 * the rights to use, copy, modify, merge, publish, distribute, sublicense,
7 * and/or sell copies of the Software, and to permit persons to whom the
8 * Software is furnished to do so, subject to the following conditions:
9 *
10 * The above copyright notice and this permission notice (including the next
11 * paragraph) shall be included in all copies or substantial portions of the
12 * Software.
13 *
14 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
15 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
16 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL
17 * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
18 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
19 * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
20 * IN THE SOFTWARE.
21 */
22
23#include "nir.h"
24
25const nir_intrinsic_info nir_intrinsic_infos[nir_num_intrinsics] = {
26{
27   .name = "accept_ray_intersection",
28   .num_srcs = 0,
29   .has_dest = false,
30   .dest_components = 0,
31   .dest_bit_sizes = 0x0,
32   .bit_size_src = -1,
33   .num_indices = 0,
34   .flags = 0,
35},
36{
37   .name = "addr_mode_is",
38   .num_srcs = 1,
39   .src_components = {
40      -1
41   },
42   .has_dest = true,
43   .dest_components = 1,
44   .dest_bit_sizes = 0x0,
45   .bit_size_src = -1,
46   .num_indices = 1,
47   .indices = {
48      NIR_INTRINSIC_MEMORY_MODES,
49   },
50   .index_map = {
51      [NIR_INTRINSIC_MEMORY_MODES] = 1,
52    },
53   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
54},
55{
56   .name = "alloc_vertices_and_primitives_amd",
57   .num_srcs = 2,
58   .src_components = {
59      1, 1
60   },
61   .has_dest = false,
62   .dest_components = 0,
63   .dest_bit_sizes = 0x0,
64   .bit_size_src = -1,
65   .num_indices = 0,
66   .flags = 0,
67},
68{
69   .name = "atomic_counter_add",
70   .num_srcs = 2,
71   .src_components = {
72      1, 1
73   },
74   .has_dest = true,
75   .dest_components = 1,
76   .dest_bit_sizes = 0x0,
77   .bit_size_src = -1,
78   .num_indices = 1,
79   .indices = {
80      NIR_INTRINSIC_BASE,
81   },
82   .index_map = {
83      [NIR_INTRINSIC_BASE] = 1,
84    },
85   .flags = 0,
86},
87{
88   .name = "atomic_counter_add_deref",
89   .num_srcs = 2,
90   .src_components = {
91      -1, 1
92   },
93   .has_dest = true,
94   .dest_components = 1,
95   .dest_bit_sizes = 0x0,
96   .bit_size_src = -1,
97   .num_indices = 0,
98   .flags = 0,
99},
100{
101   .name = "atomic_counter_and",
102   .num_srcs = 2,
103   .src_components = {
104      1, 1
105   },
106   .has_dest = true,
107   .dest_components = 1,
108   .dest_bit_sizes = 0x0,
109   .bit_size_src = -1,
110   .num_indices = 1,
111   .indices = {
112      NIR_INTRINSIC_BASE,
113   },
114   .index_map = {
115      [NIR_INTRINSIC_BASE] = 1,
116    },
117   .flags = 0,
118},
119{
120   .name = "atomic_counter_and_deref",
121   .num_srcs = 2,
122   .src_components = {
123      -1, 1
124   },
125   .has_dest = true,
126   .dest_components = 1,
127   .dest_bit_sizes = 0x0,
128   .bit_size_src = -1,
129   .num_indices = 0,
130   .flags = 0,
131},
132{
133   .name = "atomic_counter_comp_swap",
134   .num_srcs = 3,
135   .src_components = {
136      1, 1, 1
137   },
138   .has_dest = true,
139   .dest_components = 1,
140   .dest_bit_sizes = 0x0,
141   .bit_size_src = -1,
142   .num_indices = 1,
143   .indices = {
144      NIR_INTRINSIC_BASE,
145   },
146   .index_map = {
147      [NIR_INTRINSIC_BASE] = 1,
148    },
149   .flags = 0,
150},
151{
152   .name = "atomic_counter_comp_swap_deref",
153   .num_srcs = 3,
154   .src_components = {
155      -1, 1, 1
156   },
157   .has_dest = true,
158   .dest_components = 1,
159   .dest_bit_sizes = 0x0,
160   .bit_size_src = -1,
161   .num_indices = 0,
162   .flags = 0,
163},
164{
165   .name = "atomic_counter_exchange",
166   .num_srcs = 2,
167   .src_components = {
168      1, 1
169   },
170   .has_dest = true,
171   .dest_components = 1,
172   .dest_bit_sizes = 0x0,
173   .bit_size_src = -1,
174   .num_indices = 1,
175   .indices = {
176      NIR_INTRINSIC_BASE,
177   },
178   .index_map = {
179      [NIR_INTRINSIC_BASE] = 1,
180    },
181   .flags = 0,
182},
183{
184   .name = "atomic_counter_exchange_deref",
185   .num_srcs = 2,
186   .src_components = {
187      -1, 1
188   },
189   .has_dest = true,
190   .dest_components = 1,
191   .dest_bit_sizes = 0x0,
192   .bit_size_src = -1,
193   .num_indices = 0,
194   .flags = 0,
195},
196{
197   .name = "atomic_counter_inc",
198   .num_srcs = 1,
199   .src_components = {
200      1
201   },
202   .has_dest = true,
203   .dest_components = 1,
204   .dest_bit_sizes = 0x0,
205   .bit_size_src = -1,
206   .num_indices = 1,
207   .indices = {
208      NIR_INTRINSIC_BASE,
209   },
210   .index_map = {
211      [NIR_INTRINSIC_BASE] = 1,
212    },
213   .flags = 0,
214},
215{
216   .name = "atomic_counter_inc_deref",
217   .num_srcs = 1,
218   .src_components = {
219      -1
220   },
221   .has_dest = true,
222   .dest_components = 1,
223   .dest_bit_sizes = 0x0,
224   .bit_size_src = -1,
225   .num_indices = 0,
226   .flags = 0,
227},
228{
229   .name = "atomic_counter_max",
230   .num_srcs = 2,
231   .src_components = {
232      1, 1
233   },
234   .has_dest = true,
235   .dest_components = 1,
236   .dest_bit_sizes = 0x0,
237   .bit_size_src = -1,
238   .num_indices = 1,
239   .indices = {
240      NIR_INTRINSIC_BASE,
241   },
242   .index_map = {
243      [NIR_INTRINSIC_BASE] = 1,
244    },
245   .flags = 0,
246},
247{
248   .name = "atomic_counter_max_deref",
249   .num_srcs = 2,
250   .src_components = {
251      -1, 1
252   },
253   .has_dest = true,
254   .dest_components = 1,
255   .dest_bit_sizes = 0x0,
256   .bit_size_src = -1,
257   .num_indices = 0,
258   .flags = 0,
259},
260{
261   .name = "atomic_counter_min",
262   .num_srcs = 2,
263   .src_components = {
264      1, 1
265   },
266   .has_dest = true,
267   .dest_components = 1,
268   .dest_bit_sizes = 0x0,
269   .bit_size_src = -1,
270   .num_indices = 1,
271   .indices = {
272      NIR_INTRINSIC_BASE,
273   },
274   .index_map = {
275      [NIR_INTRINSIC_BASE] = 1,
276    },
277   .flags = 0,
278},
279{
280   .name = "atomic_counter_min_deref",
281   .num_srcs = 2,
282   .src_components = {
283      -1, 1
284   },
285   .has_dest = true,
286   .dest_components = 1,
287   .dest_bit_sizes = 0x0,
288   .bit_size_src = -1,
289   .num_indices = 0,
290   .flags = 0,
291},
292{
293   .name = "atomic_counter_or",
294   .num_srcs = 2,
295   .src_components = {
296      1, 1
297   },
298   .has_dest = true,
299   .dest_components = 1,
300   .dest_bit_sizes = 0x0,
301   .bit_size_src = -1,
302   .num_indices = 1,
303   .indices = {
304      NIR_INTRINSIC_BASE,
305   },
306   .index_map = {
307      [NIR_INTRINSIC_BASE] = 1,
308    },
309   .flags = 0,
310},
311{
312   .name = "atomic_counter_or_deref",
313   .num_srcs = 2,
314   .src_components = {
315      -1, 1
316   },
317   .has_dest = true,
318   .dest_components = 1,
319   .dest_bit_sizes = 0x0,
320   .bit_size_src = -1,
321   .num_indices = 0,
322   .flags = 0,
323},
324{
325   .name = "atomic_counter_post_dec",
326   .num_srcs = 1,
327   .src_components = {
328      1
329   },
330   .has_dest = true,
331   .dest_components = 1,
332   .dest_bit_sizes = 0x0,
333   .bit_size_src = -1,
334   .num_indices = 1,
335   .indices = {
336      NIR_INTRINSIC_BASE,
337   },
338   .index_map = {
339      [NIR_INTRINSIC_BASE] = 1,
340    },
341   .flags = 0,
342},
343{
344   .name = "atomic_counter_post_dec_deref",
345   .num_srcs = 1,
346   .src_components = {
347      -1
348   },
349   .has_dest = true,
350   .dest_components = 1,
351   .dest_bit_sizes = 0x0,
352   .bit_size_src = -1,
353   .num_indices = 0,
354   .flags = 0,
355},
356{
357   .name = "atomic_counter_pre_dec",
358   .num_srcs = 1,
359   .src_components = {
360      1
361   },
362   .has_dest = true,
363   .dest_components = 1,
364   .dest_bit_sizes = 0x0,
365   .bit_size_src = -1,
366   .num_indices = 1,
367   .indices = {
368      NIR_INTRINSIC_BASE,
369   },
370   .index_map = {
371      [NIR_INTRINSIC_BASE] = 1,
372    },
373   .flags = 0,
374},
375{
376   .name = "atomic_counter_pre_dec_deref",
377   .num_srcs = 1,
378   .src_components = {
379      -1
380   },
381   .has_dest = true,
382   .dest_components = 1,
383   .dest_bit_sizes = 0x0,
384   .bit_size_src = -1,
385   .num_indices = 0,
386   .flags = 0,
387},
388{
389   .name = "atomic_counter_read",
390   .num_srcs = 1,
391   .src_components = {
392      1
393   },
394   .has_dest = true,
395   .dest_components = 1,
396   .dest_bit_sizes = 0x0,
397   .bit_size_src = -1,
398   .num_indices = 1,
399   .indices = {
400      NIR_INTRINSIC_BASE,
401   },
402   .index_map = {
403      [NIR_INTRINSIC_BASE] = 1,
404    },
405   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
406},
407{
408   .name = "atomic_counter_read_deref",
409   .num_srcs = 1,
410   .src_components = {
411      -1
412   },
413   .has_dest = true,
414   .dest_components = 1,
415   .dest_bit_sizes = 0x0,
416   .bit_size_src = -1,
417   .num_indices = 0,
418   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
419},
420{
421   .name = "atomic_counter_xor",
422   .num_srcs = 2,
423   .src_components = {
424      1, 1
425   },
426   .has_dest = true,
427   .dest_components = 1,
428   .dest_bit_sizes = 0x0,
429   .bit_size_src = -1,
430   .num_indices = 1,
431   .indices = {
432      NIR_INTRINSIC_BASE,
433   },
434   .index_map = {
435      [NIR_INTRINSIC_BASE] = 1,
436    },
437   .flags = 0,
438},
439{
440   .name = "atomic_counter_xor_deref",
441   .num_srcs = 2,
442   .src_components = {
443      -1, 1
444   },
445   .has_dest = true,
446   .dest_components = 1,
447   .dest_bit_sizes = 0x0,
448   .bit_size_src = -1,
449   .num_indices = 0,
450   .flags = 0,
451},
452{
453   .name = "ballot",
454   .num_srcs = 1,
455   .src_components = {
456      1
457   },
458   .has_dest = true,
459   .dest_components = 0,
460   .dest_bit_sizes = 0x0,
461   .bit_size_src = -1,
462   .num_indices = 0,
463   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
464},
465{
466   .name = "ballot_bit_count_exclusive",
467   .num_srcs = 1,
468   .src_components = {
469      4
470   },
471   .has_dest = true,
472   .dest_components = 1,
473   .dest_bit_sizes = 0x0,
474   .bit_size_src = -1,
475   .num_indices = 0,
476   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
477},
478{
479   .name = "ballot_bit_count_inclusive",
480   .num_srcs = 1,
481   .src_components = {
482      4
483   },
484   .has_dest = true,
485   .dest_components = 1,
486   .dest_bit_sizes = 0x0,
487   .bit_size_src = -1,
488   .num_indices = 0,
489   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
490},
491{
492   .name = "ballot_bit_count_reduce",
493   .num_srcs = 1,
494   .src_components = {
495      4
496   },
497   .has_dest = true,
498   .dest_components = 1,
499   .dest_bit_sizes = 0x0,
500   .bit_size_src = -1,
501   .num_indices = 0,
502   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
503},
504{
505   .name = "ballot_bitfield_extract",
506   .num_srcs = 2,
507   .src_components = {
508      4, 1
509   },
510   .has_dest = true,
511   .dest_components = 1,
512   .dest_bit_sizes = 0x0,
513   .bit_size_src = -1,
514   .num_indices = 0,
515   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
516},
517{
518   .name = "ballot_find_lsb",
519   .num_srcs = 1,
520   .src_components = {
521      4
522   },
523   .has_dest = true,
524   .dest_components = 1,
525   .dest_bit_sizes = 0x0,
526   .bit_size_src = -1,
527   .num_indices = 0,
528   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
529},
530{
531   .name = "ballot_find_msb",
532   .num_srcs = 1,
533   .src_components = {
534      4
535   },
536   .has_dest = true,
537   .dest_components = 1,
538   .dest_bit_sizes = 0x0,
539   .bit_size_src = -1,
540   .num_indices = 0,
541   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
542},
543{
544   .name = "begin_invocation_interlock",
545   .num_srcs = 0,
546   .has_dest = false,
547   .dest_components = 0,
548   .dest_bit_sizes = 0x0,
549   .bit_size_src = -1,
550   .num_indices = 0,
551   .flags = 0,
552},
553{
554   .name = "bindless_image_atomic_add",
555   .num_srcs = 4,
556   .src_components = {
557      1, 4, 1, 1
558   },
559   .has_dest = true,
560   .dest_components = 1,
561   .dest_bit_sizes = 0x0,
562   .bit_size_src = -1,
563   .num_indices = 4,
564   .indices = {
565      NIR_INTRINSIC_IMAGE_DIM,
566      NIR_INTRINSIC_IMAGE_ARRAY,
567      NIR_INTRINSIC_FORMAT,
568      NIR_INTRINSIC_ACCESS,
569   },
570   .index_map = {
571      [NIR_INTRINSIC_IMAGE_DIM] = 1,
572      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
573      [NIR_INTRINSIC_FORMAT] = 3,
574      [NIR_INTRINSIC_ACCESS] = 4,
575    },
576   .flags = 0,
577},
578{
579   .name = "bindless_image_atomic_and",
580   .num_srcs = 4,
581   .src_components = {
582      1, 4, 1, 1
583   },
584   .has_dest = true,
585   .dest_components = 1,
586   .dest_bit_sizes = 0x0,
587   .bit_size_src = -1,
588   .num_indices = 4,
589   .indices = {
590      NIR_INTRINSIC_IMAGE_DIM,
591      NIR_INTRINSIC_IMAGE_ARRAY,
592      NIR_INTRINSIC_FORMAT,
593      NIR_INTRINSIC_ACCESS,
594   },
595   .index_map = {
596      [NIR_INTRINSIC_IMAGE_DIM] = 1,
597      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
598      [NIR_INTRINSIC_FORMAT] = 3,
599      [NIR_INTRINSIC_ACCESS] = 4,
600    },
601   .flags = 0,
602},
603{
604   .name = "bindless_image_atomic_comp_swap",
605   .num_srcs = 5,
606   .src_components = {
607      1, 4, 1, 1, 1
608   },
609   .has_dest = true,
610   .dest_components = 1,
611   .dest_bit_sizes = 0x0,
612   .bit_size_src = -1,
613   .num_indices = 4,
614   .indices = {
615      NIR_INTRINSIC_IMAGE_DIM,
616      NIR_INTRINSIC_IMAGE_ARRAY,
617      NIR_INTRINSIC_FORMAT,
618      NIR_INTRINSIC_ACCESS,
619   },
620   .index_map = {
621      [NIR_INTRINSIC_IMAGE_DIM] = 1,
622      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
623      [NIR_INTRINSIC_FORMAT] = 3,
624      [NIR_INTRINSIC_ACCESS] = 4,
625    },
626   .flags = 0,
627},
628{
629   .name = "bindless_image_atomic_dec_wrap",
630   .num_srcs = 4,
631   .src_components = {
632      1, 4, 1, 1
633   },
634   .has_dest = true,
635   .dest_components = 1,
636   .dest_bit_sizes = 0x0,
637   .bit_size_src = -1,
638   .num_indices = 4,
639   .indices = {
640      NIR_INTRINSIC_IMAGE_DIM,
641      NIR_INTRINSIC_IMAGE_ARRAY,
642      NIR_INTRINSIC_FORMAT,
643      NIR_INTRINSIC_ACCESS,
644   },
645   .index_map = {
646      [NIR_INTRINSIC_IMAGE_DIM] = 1,
647      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
648      [NIR_INTRINSIC_FORMAT] = 3,
649      [NIR_INTRINSIC_ACCESS] = 4,
650    },
651   .flags = 0,
652},
653{
654   .name = "bindless_image_atomic_exchange",
655   .num_srcs = 4,
656   .src_components = {
657      1, 4, 1, 1
658   },
659   .has_dest = true,
660   .dest_components = 1,
661   .dest_bit_sizes = 0x0,
662   .bit_size_src = -1,
663   .num_indices = 4,
664   .indices = {
665      NIR_INTRINSIC_IMAGE_DIM,
666      NIR_INTRINSIC_IMAGE_ARRAY,
667      NIR_INTRINSIC_FORMAT,
668      NIR_INTRINSIC_ACCESS,
669   },
670   .index_map = {
671      [NIR_INTRINSIC_IMAGE_DIM] = 1,
672      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
673      [NIR_INTRINSIC_FORMAT] = 3,
674      [NIR_INTRINSIC_ACCESS] = 4,
675    },
676   .flags = 0,
677},
678{
679   .name = "bindless_image_atomic_fadd",
680   .num_srcs = 4,
681   .src_components = {
682      1, 4, 1, 1
683   },
684   .has_dest = true,
685   .dest_components = 1,
686   .dest_bit_sizes = 0x0,
687   .bit_size_src = -1,
688   .num_indices = 4,
689   .indices = {
690      NIR_INTRINSIC_IMAGE_DIM,
691      NIR_INTRINSIC_IMAGE_ARRAY,
692      NIR_INTRINSIC_FORMAT,
693      NIR_INTRINSIC_ACCESS,
694   },
695   .index_map = {
696      [NIR_INTRINSIC_IMAGE_DIM] = 1,
697      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
698      [NIR_INTRINSIC_FORMAT] = 3,
699      [NIR_INTRINSIC_ACCESS] = 4,
700    },
701   .flags = 0,
702},
703{
704   .name = "bindless_image_atomic_fmax",
705   .num_srcs = 4,
706   .src_components = {
707      1, 4, 1, 1
708   },
709   .has_dest = true,
710   .dest_components = 1,
711   .dest_bit_sizes = 0x0,
712   .bit_size_src = -1,
713   .num_indices = 4,
714   .indices = {
715      NIR_INTRINSIC_IMAGE_DIM,
716      NIR_INTRINSIC_IMAGE_ARRAY,
717      NIR_INTRINSIC_FORMAT,
718      NIR_INTRINSIC_ACCESS,
719   },
720   .index_map = {
721      [NIR_INTRINSIC_IMAGE_DIM] = 1,
722      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
723      [NIR_INTRINSIC_FORMAT] = 3,
724      [NIR_INTRINSIC_ACCESS] = 4,
725    },
726   .flags = 0,
727},
728{
729   .name = "bindless_image_atomic_fmin",
730   .num_srcs = 4,
731   .src_components = {
732      1, 4, 1, 1
733   },
734   .has_dest = true,
735   .dest_components = 1,
736   .dest_bit_sizes = 0x0,
737   .bit_size_src = -1,
738   .num_indices = 4,
739   .indices = {
740      NIR_INTRINSIC_IMAGE_DIM,
741      NIR_INTRINSIC_IMAGE_ARRAY,
742      NIR_INTRINSIC_FORMAT,
743      NIR_INTRINSIC_ACCESS,
744   },
745   .index_map = {
746      [NIR_INTRINSIC_IMAGE_DIM] = 1,
747      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
748      [NIR_INTRINSIC_FORMAT] = 3,
749      [NIR_INTRINSIC_ACCESS] = 4,
750    },
751   .flags = 0,
752},
753{
754   .name = "bindless_image_atomic_imax",
755   .num_srcs = 4,
756   .src_components = {
757      1, 4, 1, 1
758   },
759   .has_dest = true,
760   .dest_components = 1,
761   .dest_bit_sizes = 0x0,
762   .bit_size_src = -1,
763   .num_indices = 4,
764   .indices = {
765      NIR_INTRINSIC_IMAGE_DIM,
766      NIR_INTRINSIC_IMAGE_ARRAY,
767      NIR_INTRINSIC_FORMAT,
768      NIR_INTRINSIC_ACCESS,
769   },
770   .index_map = {
771      [NIR_INTRINSIC_IMAGE_DIM] = 1,
772      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
773      [NIR_INTRINSIC_FORMAT] = 3,
774      [NIR_INTRINSIC_ACCESS] = 4,
775    },
776   .flags = 0,
777},
778{
779   .name = "bindless_image_atomic_imin",
780   .num_srcs = 4,
781   .src_components = {
782      1, 4, 1, 1
783   },
784   .has_dest = true,
785   .dest_components = 1,
786   .dest_bit_sizes = 0x0,
787   .bit_size_src = -1,
788   .num_indices = 4,
789   .indices = {
790      NIR_INTRINSIC_IMAGE_DIM,
791      NIR_INTRINSIC_IMAGE_ARRAY,
792      NIR_INTRINSIC_FORMAT,
793      NIR_INTRINSIC_ACCESS,
794   },
795   .index_map = {
796      [NIR_INTRINSIC_IMAGE_DIM] = 1,
797      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
798      [NIR_INTRINSIC_FORMAT] = 3,
799      [NIR_INTRINSIC_ACCESS] = 4,
800    },
801   .flags = 0,
802},
803{
804   .name = "bindless_image_atomic_inc_wrap",
805   .num_srcs = 4,
806   .src_components = {
807      1, 4, 1, 1
808   },
809   .has_dest = true,
810   .dest_components = 1,
811   .dest_bit_sizes = 0x0,
812   .bit_size_src = -1,
813   .num_indices = 4,
814   .indices = {
815      NIR_INTRINSIC_IMAGE_DIM,
816      NIR_INTRINSIC_IMAGE_ARRAY,
817      NIR_INTRINSIC_FORMAT,
818      NIR_INTRINSIC_ACCESS,
819   },
820   .index_map = {
821      [NIR_INTRINSIC_IMAGE_DIM] = 1,
822      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
823      [NIR_INTRINSIC_FORMAT] = 3,
824      [NIR_INTRINSIC_ACCESS] = 4,
825    },
826   .flags = 0,
827},
828{
829   .name = "bindless_image_atomic_or",
830   .num_srcs = 4,
831   .src_components = {
832      1, 4, 1, 1
833   },
834   .has_dest = true,
835   .dest_components = 1,
836   .dest_bit_sizes = 0x0,
837   .bit_size_src = -1,
838   .num_indices = 4,
839   .indices = {
840      NIR_INTRINSIC_IMAGE_DIM,
841      NIR_INTRINSIC_IMAGE_ARRAY,
842      NIR_INTRINSIC_FORMAT,
843      NIR_INTRINSIC_ACCESS,
844   },
845   .index_map = {
846      [NIR_INTRINSIC_IMAGE_DIM] = 1,
847      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
848      [NIR_INTRINSIC_FORMAT] = 3,
849      [NIR_INTRINSIC_ACCESS] = 4,
850    },
851   .flags = 0,
852},
853{
854   .name = "bindless_image_atomic_umax",
855   .num_srcs = 4,
856   .src_components = {
857      1, 4, 1, 1
858   },
859   .has_dest = true,
860   .dest_components = 1,
861   .dest_bit_sizes = 0x0,
862   .bit_size_src = -1,
863   .num_indices = 4,
864   .indices = {
865      NIR_INTRINSIC_IMAGE_DIM,
866      NIR_INTRINSIC_IMAGE_ARRAY,
867      NIR_INTRINSIC_FORMAT,
868      NIR_INTRINSIC_ACCESS,
869   },
870   .index_map = {
871      [NIR_INTRINSIC_IMAGE_DIM] = 1,
872      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
873      [NIR_INTRINSIC_FORMAT] = 3,
874      [NIR_INTRINSIC_ACCESS] = 4,
875    },
876   .flags = 0,
877},
878{
879   .name = "bindless_image_atomic_umin",
880   .num_srcs = 4,
881   .src_components = {
882      1, 4, 1, 1
883   },
884   .has_dest = true,
885   .dest_components = 1,
886   .dest_bit_sizes = 0x0,
887   .bit_size_src = -1,
888   .num_indices = 4,
889   .indices = {
890      NIR_INTRINSIC_IMAGE_DIM,
891      NIR_INTRINSIC_IMAGE_ARRAY,
892      NIR_INTRINSIC_FORMAT,
893      NIR_INTRINSIC_ACCESS,
894   },
895   .index_map = {
896      [NIR_INTRINSIC_IMAGE_DIM] = 1,
897      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
898      [NIR_INTRINSIC_FORMAT] = 3,
899      [NIR_INTRINSIC_ACCESS] = 4,
900    },
901   .flags = 0,
902},
903{
904   .name = "bindless_image_atomic_xor",
905   .num_srcs = 4,
906   .src_components = {
907      1, 4, 1, 1
908   },
909   .has_dest = true,
910   .dest_components = 1,
911   .dest_bit_sizes = 0x0,
912   .bit_size_src = -1,
913   .num_indices = 4,
914   .indices = {
915      NIR_INTRINSIC_IMAGE_DIM,
916      NIR_INTRINSIC_IMAGE_ARRAY,
917      NIR_INTRINSIC_FORMAT,
918      NIR_INTRINSIC_ACCESS,
919   },
920   .index_map = {
921      [NIR_INTRINSIC_IMAGE_DIM] = 1,
922      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
923      [NIR_INTRINSIC_FORMAT] = 3,
924      [NIR_INTRINSIC_ACCESS] = 4,
925    },
926   .flags = 0,
927},
928{
929   .name = "bindless_image_format",
930   .num_srcs = 1,
931   .src_components = {
932      1
933   },
934   .has_dest = true,
935   .dest_components = 1,
936   .dest_bit_sizes = 0x0,
937   .bit_size_src = -1,
938   .num_indices = 4,
939   .indices = {
940      NIR_INTRINSIC_IMAGE_DIM,
941      NIR_INTRINSIC_IMAGE_ARRAY,
942      NIR_INTRINSIC_FORMAT,
943      NIR_INTRINSIC_ACCESS,
944   },
945   .index_map = {
946      [NIR_INTRINSIC_IMAGE_DIM] = 1,
947      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
948      [NIR_INTRINSIC_FORMAT] = 3,
949      [NIR_INTRINSIC_ACCESS] = 4,
950    },
951   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
952},
953{
954   .name = "bindless_image_load",
955   .num_srcs = 4,
956   .src_components = {
957      1, 4, 1, 1
958   },
959   .has_dest = true,
960   .dest_components = 0,
961   .dest_bit_sizes = 0x0,
962   .bit_size_src = -1,
963   .num_indices = 5,
964   .indices = {
965      NIR_INTRINSIC_IMAGE_DIM,
966      NIR_INTRINSIC_IMAGE_ARRAY,
967      NIR_INTRINSIC_FORMAT,
968      NIR_INTRINSIC_ACCESS,
969      NIR_INTRINSIC_DEST_TYPE,
970   },
971   .index_map = {
972      [NIR_INTRINSIC_IMAGE_DIM] = 1,
973      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
974      [NIR_INTRINSIC_FORMAT] = 3,
975      [NIR_INTRINSIC_ACCESS] = 4,
976      [NIR_INTRINSIC_DEST_TYPE] = 5,
977    },
978   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
979},
980{
981   .name = "bindless_image_load_raw_intel",
982   .num_srcs = 2,
983   .src_components = {
984      1, 1
985   },
986   .has_dest = true,
987   .dest_components = 0,
988   .dest_bit_sizes = 0x0,
989   .bit_size_src = -1,
990   .num_indices = 4,
991   .indices = {
992      NIR_INTRINSIC_IMAGE_DIM,
993      NIR_INTRINSIC_IMAGE_ARRAY,
994      NIR_INTRINSIC_FORMAT,
995      NIR_INTRINSIC_ACCESS,
996   },
997   .index_map = {
998      [NIR_INTRINSIC_IMAGE_DIM] = 1,
999      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
1000      [NIR_INTRINSIC_FORMAT] = 3,
1001      [NIR_INTRINSIC_ACCESS] = 4,
1002    },
1003   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
1004},
1005{
1006   .name = "bindless_image_order",
1007   .num_srcs = 1,
1008   .src_components = {
1009      1
1010   },
1011   .has_dest = true,
1012   .dest_components = 1,
1013   .dest_bit_sizes = 0x0,
1014   .bit_size_src = -1,
1015   .num_indices = 4,
1016   .indices = {
1017      NIR_INTRINSIC_IMAGE_DIM,
1018      NIR_INTRINSIC_IMAGE_ARRAY,
1019      NIR_INTRINSIC_FORMAT,
1020      NIR_INTRINSIC_ACCESS,
1021   },
1022   .index_map = {
1023      [NIR_INTRINSIC_IMAGE_DIM] = 1,
1024      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
1025      [NIR_INTRINSIC_FORMAT] = 3,
1026      [NIR_INTRINSIC_ACCESS] = 4,
1027    },
1028   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
1029},
1030{
1031   .name = "bindless_image_samples",
1032   .num_srcs = 1,
1033   .src_components = {
1034      1
1035   },
1036   .has_dest = true,
1037   .dest_components = 1,
1038   .dest_bit_sizes = 0x0,
1039   .bit_size_src = -1,
1040   .num_indices = 4,
1041   .indices = {
1042      NIR_INTRINSIC_IMAGE_DIM,
1043      NIR_INTRINSIC_IMAGE_ARRAY,
1044      NIR_INTRINSIC_FORMAT,
1045      NIR_INTRINSIC_ACCESS,
1046   },
1047   .index_map = {
1048      [NIR_INTRINSIC_IMAGE_DIM] = 1,
1049      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
1050      [NIR_INTRINSIC_FORMAT] = 3,
1051      [NIR_INTRINSIC_ACCESS] = 4,
1052    },
1053   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
1054},
1055{
1056   .name = "bindless_image_size",
1057   .num_srcs = 2,
1058   .src_components = {
1059      1, 1
1060   },
1061   .has_dest = true,
1062   .dest_components = 0,
1063   .dest_bit_sizes = 0x0,
1064   .bit_size_src = -1,
1065   .num_indices = 4,
1066   .indices = {
1067      NIR_INTRINSIC_IMAGE_DIM,
1068      NIR_INTRINSIC_IMAGE_ARRAY,
1069      NIR_INTRINSIC_FORMAT,
1070      NIR_INTRINSIC_ACCESS,
1071   },
1072   .index_map = {
1073      [NIR_INTRINSIC_IMAGE_DIM] = 1,
1074      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
1075      [NIR_INTRINSIC_FORMAT] = 3,
1076      [NIR_INTRINSIC_ACCESS] = 4,
1077    },
1078   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
1079},
1080{
1081   .name = "bindless_image_sparse_load",
1082   .num_srcs = 4,
1083   .src_components = {
1084      1, 4, 1, 1
1085   },
1086   .has_dest = true,
1087   .dest_components = 0,
1088   .dest_bit_sizes = 0x0,
1089   .bit_size_src = -1,
1090   .num_indices = 5,
1091   .indices = {
1092      NIR_INTRINSIC_IMAGE_DIM,
1093      NIR_INTRINSIC_IMAGE_ARRAY,
1094      NIR_INTRINSIC_FORMAT,
1095      NIR_INTRINSIC_ACCESS,
1096      NIR_INTRINSIC_DEST_TYPE,
1097   },
1098   .index_map = {
1099      [NIR_INTRINSIC_IMAGE_DIM] = 1,
1100      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
1101      [NIR_INTRINSIC_FORMAT] = 3,
1102      [NIR_INTRINSIC_ACCESS] = 4,
1103      [NIR_INTRINSIC_DEST_TYPE] = 5,
1104    },
1105   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
1106},
1107{
1108   .name = "bindless_image_store",
1109   .num_srcs = 5,
1110   .src_components = {
1111      1, 4, 1, 0, 1
1112   },
1113   .has_dest = false,
1114   .dest_components = 0,
1115   .dest_bit_sizes = 0x0,
1116   .bit_size_src = -1,
1117   .num_indices = 5,
1118   .indices = {
1119      NIR_INTRINSIC_IMAGE_DIM,
1120      NIR_INTRINSIC_IMAGE_ARRAY,
1121      NIR_INTRINSIC_FORMAT,
1122      NIR_INTRINSIC_ACCESS,
1123      NIR_INTRINSIC_SRC_TYPE,
1124   },
1125   .index_map = {
1126      [NIR_INTRINSIC_IMAGE_DIM] = 1,
1127      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
1128      [NIR_INTRINSIC_FORMAT] = 3,
1129      [NIR_INTRINSIC_ACCESS] = 4,
1130      [NIR_INTRINSIC_SRC_TYPE] = 5,
1131    },
1132   .flags = 0,
1133},
1134{
1135   .name = "bindless_image_store_raw_intel",
1136   .num_srcs = 3,
1137   .src_components = {
1138      1, 1, 0
1139   },
1140   .has_dest = false,
1141   .dest_components = 0,
1142   .dest_bit_sizes = 0x0,
1143   .bit_size_src = -1,
1144   .num_indices = 4,
1145   .indices = {
1146      NIR_INTRINSIC_IMAGE_DIM,
1147      NIR_INTRINSIC_IMAGE_ARRAY,
1148      NIR_INTRINSIC_FORMAT,
1149      NIR_INTRINSIC_ACCESS,
1150   },
1151   .index_map = {
1152      [NIR_INTRINSIC_IMAGE_DIM] = 1,
1153      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
1154      [NIR_INTRINSIC_FORMAT] = 3,
1155      [NIR_INTRINSIC_ACCESS] = 4,
1156    },
1157   .flags = 0,
1158},
1159{
1160   .name = "bindless_resource_ir3",
1161   .num_srcs = 1,
1162   .src_components = {
1163      1
1164   },
1165   .has_dest = true,
1166   .dest_components = 1,
1167   .dest_bit_sizes = 0x0,
1168   .bit_size_src = -1,
1169   .num_indices = 1,
1170   .indices = {
1171      NIR_INTRINSIC_DESC_SET,
1172   },
1173   .index_map = {
1174      [NIR_INTRINSIC_DESC_SET] = 1,
1175    },
1176   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
1177},
1178{
1179   .name = "btd_retire_intel",
1180   .num_srcs = 0,
1181   .has_dest = false,
1182   .dest_components = 0,
1183   .dest_bit_sizes = 0x0,
1184   .bit_size_src = -1,
1185   .num_indices = 0,
1186   .flags = 0,
1187},
1188{
1189   .name = "btd_spawn_intel",
1190   .num_srcs = 2,
1191   .src_components = {
1192      1, 1
1193   },
1194   .has_dest = false,
1195   .dest_components = 0,
1196   .dest_bit_sizes = 0x0,
1197   .bit_size_src = -1,
1198   .num_indices = 0,
1199   .flags = 0,
1200},
1201{
1202   .name = "btd_stack_push_intel",
1203   .num_srcs = 0,
1204   .has_dest = false,
1205   .dest_components = 0,
1206   .dest_bit_sizes = 0x0,
1207   .bit_size_src = -1,
1208   .num_indices = 1,
1209   .indices = {
1210      NIR_INTRINSIC_STACK_SIZE,
1211   },
1212   .index_map = {
1213      [NIR_INTRINSIC_STACK_SIZE] = 1,
1214    },
1215   .flags = 0,
1216},
1217{
1218   .name = "bvh64_intersect_ray_amd",
1219   .num_srcs = 6,
1220   .src_components = {
1221      4, 2, 1, 3, 3, 3
1222   },
1223   .has_dest = true,
1224   .dest_components = 4,
1225   .dest_bit_sizes = 0x0,
1226   .bit_size_src = -1,
1227   .num_indices = 0,
1228   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
1229},
1230{
1231   .name = "byte_permute_amd",
1232   .num_srcs = 3,
1233   .src_components = {
1234      1, 1, 1
1235   },
1236   .has_dest = true,
1237   .dest_components = 1,
1238   .dest_bit_sizes = 0x20,
1239   .bit_size_src = -1,
1240   .num_indices = 0,
1241   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
1242},
1243{
1244   .name = "cond_end_ir3",
1245   .num_srcs = 1,
1246   .src_components = {
1247      1
1248   },
1249   .has_dest = false,
1250   .dest_components = 0,
1251   .dest_bit_sizes = 0x0,
1252   .bit_size_src = -1,
1253   .num_indices = 0,
1254   .flags = 0,
1255},
1256{
1257   .name = "control_barrier",
1258   .num_srcs = 0,
1259   .has_dest = false,
1260   .dest_components = 0,
1261   .dest_bit_sizes = 0x0,
1262   .bit_size_src = -1,
1263   .num_indices = 0,
1264   .flags = 0,
1265},
1266{
1267   .name = "convert_alu_types",
1268   .num_srcs = 1,
1269   .src_components = {
1270      0
1271   },
1272   .has_dest = true,
1273   .dest_components = 0,
1274   .dest_bit_sizes = 0x0,
1275   .bit_size_src = -1,
1276   .num_indices = 4,
1277   .indices = {
1278      NIR_INTRINSIC_SRC_TYPE,
1279      NIR_INTRINSIC_DEST_TYPE,
1280      NIR_INTRINSIC_ROUNDING_MODE,
1281      NIR_INTRINSIC_SATURATE,
1282   },
1283   .index_map = {
1284      [NIR_INTRINSIC_SRC_TYPE] = 1,
1285      [NIR_INTRINSIC_DEST_TYPE] = 2,
1286      [NIR_INTRINSIC_ROUNDING_MODE] = 3,
1287      [NIR_INTRINSIC_SATURATE] = 4,
1288    },
1289   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
1290},
1291{
1292   .name = "copy_deref",
1293   .num_srcs = 2,
1294   .src_components = {
1295      -1, -1
1296   },
1297   .has_dest = false,
1298   .dest_components = 0,
1299   .dest_bit_sizes = 0x0,
1300   .bit_size_src = -1,
1301   .num_indices = 2,
1302   .indices = {
1303      NIR_INTRINSIC_DST_ACCESS,
1304      NIR_INTRINSIC_SRC_ACCESS,
1305   },
1306   .index_map = {
1307      [NIR_INTRINSIC_DST_ACCESS] = 1,
1308      [NIR_INTRINSIC_SRC_ACCESS] = 2,
1309    },
1310   .flags = 0,
1311},
1312{
1313   .name = "demote",
1314   .num_srcs = 0,
1315   .has_dest = false,
1316   .dest_components = 0,
1317   .dest_bit_sizes = 0x0,
1318   .bit_size_src = -1,
1319   .num_indices = 0,
1320   .flags = 0,
1321},
1322{
1323   .name = "demote_if",
1324   .num_srcs = 1,
1325   .src_components = {
1326      1
1327   },
1328   .has_dest = false,
1329   .dest_components = 0,
1330   .dest_bit_sizes = 0x0,
1331   .bit_size_src = -1,
1332   .num_indices = 0,
1333   .flags = 0,
1334},
1335{
1336   .name = "deref_atomic_add",
1337   .num_srcs = 2,
1338   .src_components = {
1339      -1, 1
1340   },
1341   .has_dest = true,
1342   .dest_components = 1,
1343   .dest_bit_sizes = 0x0,
1344   .bit_size_src = -1,
1345   .num_indices = 1,
1346   .indices = {
1347      NIR_INTRINSIC_ACCESS,
1348   },
1349   .index_map = {
1350      [NIR_INTRINSIC_ACCESS] = 1,
1351    },
1352   .flags = 0,
1353},
1354{
1355   .name = "deref_atomic_and",
1356   .num_srcs = 2,
1357   .src_components = {
1358      -1, 1
1359   },
1360   .has_dest = true,
1361   .dest_components = 1,
1362   .dest_bit_sizes = 0x0,
1363   .bit_size_src = -1,
1364   .num_indices = 1,
1365   .indices = {
1366      NIR_INTRINSIC_ACCESS,
1367   },
1368   .index_map = {
1369      [NIR_INTRINSIC_ACCESS] = 1,
1370    },
1371   .flags = 0,
1372},
1373{
1374   .name = "deref_atomic_comp_swap",
1375   .num_srcs = 3,
1376   .src_components = {
1377      -1, 1, 1
1378   },
1379   .has_dest = true,
1380   .dest_components = 1,
1381   .dest_bit_sizes = 0x0,
1382   .bit_size_src = -1,
1383   .num_indices = 1,
1384   .indices = {
1385      NIR_INTRINSIC_ACCESS,
1386   },
1387   .index_map = {
1388      [NIR_INTRINSIC_ACCESS] = 1,
1389    },
1390   .flags = 0,
1391},
1392{
1393   .name = "deref_atomic_exchange",
1394   .num_srcs = 2,
1395   .src_components = {
1396      -1, 1
1397   },
1398   .has_dest = true,
1399   .dest_components = 1,
1400   .dest_bit_sizes = 0x0,
1401   .bit_size_src = -1,
1402   .num_indices = 1,
1403   .indices = {
1404      NIR_INTRINSIC_ACCESS,
1405   },
1406   .index_map = {
1407      [NIR_INTRINSIC_ACCESS] = 1,
1408    },
1409   .flags = 0,
1410},
1411{
1412   .name = "deref_atomic_fadd",
1413   .num_srcs = 2,
1414   .src_components = {
1415      -1, 1
1416   },
1417   .has_dest = true,
1418   .dest_components = 1,
1419   .dest_bit_sizes = 0x0,
1420   .bit_size_src = -1,
1421   .num_indices = 1,
1422   .indices = {
1423      NIR_INTRINSIC_ACCESS,
1424   },
1425   .index_map = {
1426      [NIR_INTRINSIC_ACCESS] = 1,
1427    },
1428   .flags = 0,
1429},
1430{
1431   .name = "deref_atomic_fcomp_swap",
1432   .num_srcs = 3,
1433   .src_components = {
1434      -1, 1, 1
1435   },
1436   .has_dest = true,
1437   .dest_components = 1,
1438   .dest_bit_sizes = 0x0,
1439   .bit_size_src = -1,
1440   .num_indices = 1,
1441   .indices = {
1442      NIR_INTRINSIC_ACCESS,
1443   },
1444   .index_map = {
1445      [NIR_INTRINSIC_ACCESS] = 1,
1446    },
1447   .flags = 0,
1448},
1449{
1450   .name = "deref_atomic_fmax",
1451   .num_srcs = 2,
1452   .src_components = {
1453      -1, 1
1454   },
1455   .has_dest = true,
1456   .dest_components = 1,
1457   .dest_bit_sizes = 0x0,
1458   .bit_size_src = -1,
1459   .num_indices = 1,
1460   .indices = {
1461      NIR_INTRINSIC_ACCESS,
1462   },
1463   .index_map = {
1464      [NIR_INTRINSIC_ACCESS] = 1,
1465    },
1466   .flags = 0,
1467},
1468{
1469   .name = "deref_atomic_fmin",
1470   .num_srcs = 2,
1471   .src_components = {
1472      -1, 1
1473   },
1474   .has_dest = true,
1475   .dest_components = 1,
1476   .dest_bit_sizes = 0x0,
1477   .bit_size_src = -1,
1478   .num_indices = 1,
1479   .indices = {
1480      NIR_INTRINSIC_ACCESS,
1481   },
1482   .index_map = {
1483      [NIR_INTRINSIC_ACCESS] = 1,
1484    },
1485   .flags = 0,
1486},
1487{
1488   .name = "deref_atomic_imax",
1489   .num_srcs = 2,
1490   .src_components = {
1491      -1, 1
1492   },
1493   .has_dest = true,
1494   .dest_components = 1,
1495   .dest_bit_sizes = 0x0,
1496   .bit_size_src = -1,
1497   .num_indices = 1,
1498   .indices = {
1499      NIR_INTRINSIC_ACCESS,
1500   },
1501   .index_map = {
1502      [NIR_INTRINSIC_ACCESS] = 1,
1503    },
1504   .flags = 0,
1505},
1506{
1507   .name = "deref_atomic_imin",
1508   .num_srcs = 2,
1509   .src_components = {
1510      -1, 1
1511   },
1512   .has_dest = true,
1513   .dest_components = 1,
1514   .dest_bit_sizes = 0x0,
1515   .bit_size_src = -1,
1516   .num_indices = 1,
1517   .indices = {
1518      NIR_INTRINSIC_ACCESS,
1519   },
1520   .index_map = {
1521      [NIR_INTRINSIC_ACCESS] = 1,
1522    },
1523   .flags = 0,
1524},
1525{
1526   .name = "deref_atomic_or",
1527   .num_srcs = 2,
1528   .src_components = {
1529      -1, 1
1530   },
1531   .has_dest = true,
1532   .dest_components = 1,
1533   .dest_bit_sizes = 0x0,
1534   .bit_size_src = -1,
1535   .num_indices = 1,
1536   .indices = {
1537      NIR_INTRINSIC_ACCESS,
1538   },
1539   .index_map = {
1540      [NIR_INTRINSIC_ACCESS] = 1,
1541    },
1542   .flags = 0,
1543},
1544{
1545   .name = "deref_atomic_umax",
1546   .num_srcs = 2,
1547   .src_components = {
1548      -1, 1
1549   },
1550   .has_dest = true,
1551   .dest_components = 1,
1552   .dest_bit_sizes = 0x0,
1553   .bit_size_src = -1,
1554   .num_indices = 1,
1555   .indices = {
1556      NIR_INTRINSIC_ACCESS,
1557   },
1558   .index_map = {
1559      [NIR_INTRINSIC_ACCESS] = 1,
1560    },
1561   .flags = 0,
1562},
1563{
1564   .name = "deref_atomic_umin",
1565   .num_srcs = 2,
1566   .src_components = {
1567      -1, 1
1568   },
1569   .has_dest = true,
1570   .dest_components = 1,
1571   .dest_bit_sizes = 0x0,
1572   .bit_size_src = -1,
1573   .num_indices = 1,
1574   .indices = {
1575      NIR_INTRINSIC_ACCESS,
1576   },
1577   .index_map = {
1578      [NIR_INTRINSIC_ACCESS] = 1,
1579    },
1580   .flags = 0,
1581},
1582{
1583   .name = "deref_atomic_xor",
1584   .num_srcs = 2,
1585   .src_components = {
1586      -1, 1
1587   },
1588   .has_dest = true,
1589   .dest_components = 1,
1590   .dest_bit_sizes = 0x0,
1591   .bit_size_src = -1,
1592   .num_indices = 1,
1593   .indices = {
1594      NIR_INTRINSIC_ACCESS,
1595   },
1596   .index_map = {
1597      [NIR_INTRINSIC_ACCESS] = 1,
1598    },
1599   .flags = 0,
1600},
1601{
1602   .name = "deref_buffer_array_length",
1603   .num_srcs = 1,
1604   .src_components = {
1605      -1
1606   },
1607   .has_dest = true,
1608   .dest_components = 1,
1609   .dest_bit_sizes = 0x0,
1610   .bit_size_src = -1,
1611   .num_indices = 1,
1612   .indices = {
1613      NIR_INTRINSIC_ACCESS,
1614   },
1615   .index_map = {
1616      [NIR_INTRINSIC_ACCESS] = 1,
1617    },
1618   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
1619},
1620{
1621   .name = "deref_mode_is",
1622   .num_srcs = 1,
1623   .src_components = {
1624      -1
1625   },
1626   .has_dest = true,
1627   .dest_components = 1,
1628   .dest_bit_sizes = 0x0,
1629   .bit_size_src = -1,
1630   .num_indices = 1,
1631   .indices = {
1632      NIR_INTRINSIC_MEMORY_MODES,
1633   },
1634   .index_map = {
1635      [NIR_INTRINSIC_MEMORY_MODES] = 1,
1636    },
1637   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
1638},
1639{
1640   .name = "discard",
1641   .num_srcs = 0,
1642   .has_dest = false,
1643   .dest_components = 0,
1644   .dest_bit_sizes = 0x0,
1645   .bit_size_src = -1,
1646   .num_indices = 0,
1647   .flags = 0,
1648},
1649{
1650   .name = "discard_if",
1651   .num_srcs = 1,
1652   .src_components = {
1653      1
1654   },
1655   .has_dest = false,
1656   .dest_components = 0,
1657   .dest_bit_sizes = 0x0,
1658   .bit_size_src = -1,
1659   .num_indices = 0,
1660   .flags = 0,
1661},
1662{
1663   .name = "elect",
1664   .num_srcs = 0,
1665   .has_dest = true,
1666   .dest_components = 1,
1667   .dest_bit_sizes = 0x0,
1668   .bit_size_src = -1,
1669   .num_indices = 0,
1670   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
1671},
1672{
1673   .name = "emit_vertex",
1674   .num_srcs = 0,
1675   .has_dest = false,
1676   .dest_components = 0,
1677   .dest_bit_sizes = 0x0,
1678   .bit_size_src = -1,
1679   .num_indices = 1,
1680   .indices = {
1681      NIR_INTRINSIC_STREAM_ID,
1682   },
1683   .index_map = {
1684      [NIR_INTRINSIC_STREAM_ID] = 1,
1685    },
1686   .flags = 0,
1687},
1688{
1689   .name = "emit_vertex_with_counter",
1690   .num_srcs = 2,
1691   .src_components = {
1692      1, 1
1693   },
1694   .has_dest = false,
1695   .dest_components = 0,
1696   .dest_bit_sizes = 0x0,
1697   .bit_size_src = -1,
1698   .num_indices = 1,
1699   .indices = {
1700      NIR_INTRINSIC_STREAM_ID,
1701   },
1702   .index_map = {
1703      [NIR_INTRINSIC_STREAM_ID] = 1,
1704    },
1705   .flags = 0,
1706},
1707{
1708   .name = "end_invocation_interlock",
1709   .num_srcs = 0,
1710   .has_dest = false,
1711   .dest_components = 0,
1712   .dest_bit_sizes = 0x0,
1713   .bit_size_src = -1,
1714   .num_indices = 0,
1715   .flags = 0,
1716},
1717{
1718   .name = "end_patch_ir3",
1719   .num_srcs = 0,
1720   .has_dest = false,
1721   .dest_components = 0,
1722   .dest_bit_sizes = 0x0,
1723   .bit_size_src = -1,
1724   .num_indices = 0,
1725   .flags = 0,
1726},
1727{
1728   .name = "end_primitive",
1729   .num_srcs = 0,
1730   .has_dest = false,
1731   .dest_components = 0,
1732   .dest_bit_sizes = 0x0,
1733   .bit_size_src = -1,
1734   .num_indices = 1,
1735   .indices = {
1736      NIR_INTRINSIC_STREAM_ID,
1737   },
1738   .index_map = {
1739      [NIR_INTRINSIC_STREAM_ID] = 1,
1740    },
1741   .flags = 0,
1742},
1743{
1744   .name = "end_primitive_with_counter",
1745   .num_srcs = 2,
1746   .src_components = {
1747      1, 1
1748   },
1749   .has_dest = false,
1750   .dest_components = 0,
1751   .dest_bit_sizes = 0x0,
1752   .bit_size_src = -1,
1753   .num_indices = 1,
1754   .indices = {
1755      NIR_INTRINSIC_STREAM_ID,
1756   },
1757   .index_map = {
1758      [NIR_INTRINSIC_STREAM_ID] = 1,
1759    },
1760   .flags = 0,
1761},
1762{
1763   .name = "exclusive_scan",
1764   .num_srcs = 1,
1765   .src_components = {
1766      0
1767   },
1768   .has_dest = true,
1769   .dest_components = 0,
1770   .dest_bit_sizes = 0x0,
1771   .bit_size_src = 0,
1772   .num_indices = 1,
1773   .indices = {
1774      NIR_INTRINSIC_REDUCTION_OP,
1775   },
1776   .index_map = {
1777      [NIR_INTRINSIC_REDUCTION_OP] = 1,
1778    },
1779   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
1780},
1781{
1782   .name = "execute_callable",
1783   .num_srcs = 2,
1784   .src_components = {
1785      1, -1
1786   },
1787   .has_dest = false,
1788   .dest_components = 0,
1789   .dest_bit_sizes = 0x0,
1790   .bit_size_src = -1,
1791   .num_indices = 0,
1792   .flags = 0,
1793},
1794{
1795   .name = "export_primitive_amd",
1796   .num_srcs = 1,
1797   .src_components = {
1798      1
1799   },
1800   .has_dest = false,
1801   .dest_components = 0,
1802   .dest_bit_sizes = 0x0,
1803   .bit_size_src = -1,
1804   .num_indices = 0,
1805   .flags = 0,
1806},
1807{
1808   .name = "export_vertex_amd",
1809   .num_srcs = 0,
1810   .has_dest = false,
1811   .dest_components = 0,
1812   .dest_bit_sizes = 0x0,
1813   .bit_size_src = -1,
1814   .num_indices = 0,
1815   .flags = 0,
1816},
1817{
1818   .name = "first_invocation",
1819   .num_srcs = 0,
1820   .has_dest = true,
1821   .dest_components = 1,
1822   .dest_bit_sizes = 0x20,
1823   .bit_size_src = -1,
1824   .num_indices = 0,
1825   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
1826},
1827{
1828   .name = "gds_atomic_add_amd",
1829   .num_srcs = 3,
1830   .src_components = {
1831      1, 1, 1
1832   },
1833   .has_dest = true,
1834   .dest_components = 1,
1835   .dest_bit_sizes = 0x0,
1836   .bit_size_src = -1,
1837   .num_indices = 1,
1838   .indices = {
1839      NIR_INTRINSIC_BASE,
1840   },
1841   .index_map = {
1842      [NIR_INTRINSIC_BASE] = 1,
1843    },
1844   .flags = 0,
1845},
1846{
1847   .name = "get_ssbo_size",
1848   .num_srcs = 1,
1849   .src_components = {
1850      -1
1851   },
1852   .has_dest = true,
1853   .dest_components = 1,
1854   .dest_bit_sizes = 0x20,
1855   .bit_size_src = -1,
1856   .num_indices = 1,
1857   .indices = {
1858      NIR_INTRINSIC_ACCESS,
1859   },
1860   .index_map = {
1861      [NIR_INTRINSIC_ACCESS] = 1,
1862    },
1863   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
1864},
1865{
1866   .name = "get_ubo_size",
1867   .num_srcs = 1,
1868   .src_components = {
1869      -1
1870   },
1871   .has_dest = true,
1872   .dest_components = 1,
1873   .dest_bit_sizes = 0x0,
1874   .bit_size_src = -1,
1875   .num_indices = 0,
1876   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
1877},
1878{
1879   .name = "global_atomic_add",
1880   .num_srcs = 2,
1881   .src_components = {
1882      1, 1
1883   },
1884   .has_dest = true,
1885   .dest_components = 1,
1886   .dest_bit_sizes = 0x0,
1887   .bit_size_src = -1,
1888   .num_indices = 1,
1889   .indices = {
1890      NIR_INTRINSIC_BASE,
1891   },
1892   .index_map = {
1893      [NIR_INTRINSIC_BASE] = 1,
1894    },
1895   .flags = 0,
1896},
1897{
1898   .name = "global_atomic_and",
1899   .num_srcs = 2,
1900   .src_components = {
1901      1, 1
1902   },
1903   .has_dest = true,
1904   .dest_components = 1,
1905   .dest_bit_sizes = 0x0,
1906   .bit_size_src = -1,
1907   .num_indices = 1,
1908   .indices = {
1909      NIR_INTRINSIC_BASE,
1910   },
1911   .index_map = {
1912      [NIR_INTRINSIC_BASE] = 1,
1913    },
1914   .flags = 0,
1915},
1916{
1917   .name = "global_atomic_comp_swap",
1918   .num_srcs = 3,
1919   .src_components = {
1920      1, 1, 1
1921   },
1922   .has_dest = true,
1923   .dest_components = 1,
1924   .dest_bit_sizes = 0x0,
1925   .bit_size_src = -1,
1926   .num_indices = 1,
1927   .indices = {
1928      NIR_INTRINSIC_BASE,
1929   },
1930   .index_map = {
1931      [NIR_INTRINSIC_BASE] = 1,
1932    },
1933   .flags = 0,
1934},
1935{
1936   .name = "global_atomic_exchange",
1937   .num_srcs = 2,
1938   .src_components = {
1939      1, 1
1940   },
1941   .has_dest = true,
1942   .dest_components = 1,
1943   .dest_bit_sizes = 0x0,
1944   .bit_size_src = -1,
1945   .num_indices = 1,
1946   .indices = {
1947      NIR_INTRINSIC_BASE,
1948   },
1949   .index_map = {
1950      [NIR_INTRINSIC_BASE] = 1,
1951    },
1952   .flags = 0,
1953},
1954{
1955   .name = "global_atomic_fadd",
1956   .num_srcs = 2,
1957   .src_components = {
1958      1, 1
1959   },
1960   .has_dest = true,
1961   .dest_components = 1,
1962   .dest_bit_sizes = 0x0,
1963   .bit_size_src = -1,
1964   .num_indices = 1,
1965   .indices = {
1966      NIR_INTRINSIC_BASE,
1967   },
1968   .index_map = {
1969      [NIR_INTRINSIC_BASE] = 1,
1970    },
1971   .flags = 0,
1972},
1973{
1974   .name = "global_atomic_fcomp_swap",
1975   .num_srcs = 3,
1976   .src_components = {
1977      1, 1, 1
1978   },
1979   .has_dest = true,
1980   .dest_components = 1,
1981   .dest_bit_sizes = 0x0,
1982   .bit_size_src = -1,
1983   .num_indices = 1,
1984   .indices = {
1985      NIR_INTRINSIC_BASE,
1986   },
1987   .index_map = {
1988      [NIR_INTRINSIC_BASE] = 1,
1989    },
1990   .flags = 0,
1991},
1992{
1993   .name = "global_atomic_fmax",
1994   .num_srcs = 2,
1995   .src_components = {
1996      1, 1
1997   },
1998   .has_dest = true,
1999   .dest_components = 1,
2000   .dest_bit_sizes = 0x0,
2001   .bit_size_src = -1,
2002   .num_indices = 1,
2003   .indices = {
2004      NIR_INTRINSIC_BASE,
2005   },
2006   .index_map = {
2007      [NIR_INTRINSIC_BASE] = 1,
2008    },
2009   .flags = 0,
2010},
2011{
2012   .name = "global_atomic_fmin",
2013   .num_srcs = 2,
2014   .src_components = {
2015      1, 1
2016   },
2017   .has_dest = true,
2018   .dest_components = 1,
2019   .dest_bit_sizes = 0x0,
2020   .bit_size_src = -1,
2021   .num_indices = 1,
2022   .indices = {
2023      NIR_INTRINSIC_BASE,
2024   },
2025   .index_map = {
2026      [NIR_INTRINSIC_BASE] = 1,
2027    },
2028   .flags = 0,
2029},
2030{
2031   .name = "global_atomic_imax",
2032   .num_srcs = 2,
2033   .src_components = {
2034      1, 1
2035   },
2036   .has_dest = true,
2037   .dest_components = 1,
2038   .dest_bit_sizes = 0x0,
2039   .bit_size_src = -1,
2040   .num_indices = 1,
2041   .indices = {
2042      NIR_INTRINSIC_BASE,
2043   },
2044   .index_map = {
2045      [NIR_INTRINSIC_BASE] = 1,
2046    },
2047   .flags = 0,
2048},
2049{
2050   .name = "global_atomic_imin",
2051   .num_srcs = 2,
2052   .src_components = {
2053      1, 1
2054   },
2055   .has_dest = true,
2056   .dest_components = 1,
2057   .dest_bit_sizes = 0x0,
2058   .bit_size_src = -1,
2059   .num_indices = 1,
2060   .indices = {
2061      NIR_INTRINSIC_BASE,
2062   },
2063   .index_map = {
2064      [NIR_INTRINSIC_BASE] = 1,
2065    },
2066   .flags = 0,
2067},
2068{
2069   .name = "global_atomic_or",
2070   .num_srcs = 2,
2071   .src_components = {
2072      1, 1
2073   },
2074   .has_dest = true,
2075   .dest_components = 1,
2076   .dest_bit_sizes = 0x0,
2077   .bit_size_src = -1,
2078   .num_indices = 1,
2079   .indices = {
2080      NIR_INTRINSIC_BASE,
2081   },
2082   .index_map = {
2083      [NIR_INTRINSIC_BASE] = 1,
2084    },
2085   .flags = 0,
2086},
2087{
2088   .name = "global_atomic_umax",
2089   .num_srcs = 2,
2090   .src_components = {
2091      1, 1
2092   },
2093   .has_dest = true,
2094   .dest_components = 1,
2095   .dest_bit_sizes = 0x0,
2096   .bit_size_src = -1,
2097   .num_indices = 1,
2098   .indices = {
2099      NIR_INTRINSIC_BASE,
2100   },
2101   .index_map = {
2102      [NIR_INTRINSIC_BASE] = 1,
2103    },
2104   .flags = 0,
2105},
2106{
2107   .name = "global_atomic_umin",
2108   .num_srcs = 2,
2109   .src_components = {
2110      1, 1
2111   },
2112   .has_dest = true,
2113   .dest_components = 1,
2114   .dest_bit_sizes = 0x0,
2115   .bit_size_src = -1,
2116   .num_indices = 1,
2117   .indices = {
2118      NIR_INTRINSIC_BASE,
2119   },
2120   .index_map = {
2121      [NIR_INTRINSIC_BASE] = 1,
2122    },
2123   .flags = 0,
2124},
2125{
2126   .name = "global_atomic_xor",
2127   .num_srcs = 2,
2128   .src_components = {
2129      1, 1
2130   },
2131   .has_dest = true,
2132   .dest_components = 1,
2133   .dest_bit_sizes = 0x0,
2134   .bit_size_src = -1,
2135   .num_indices = 1,
2136   .indices = {
2137      NIR_INTRINSIC_BASE,
2138   },
2139   .index_map = {
2140      [NIR_INTRINSIC_BASE] = 1,
2141    },
2142   .flags = 0,
2143},
2144{
2145   .name = "group_memory_barrier",
2146   .num_srcs = 0,
2147   .has_dest = false,
2148   .dest_components = 0,
2149   .dest_bit_sizes = 0x0,
2150   .bit_size_src = -1,
2151   .num_indices = 0,
2152   .flags = 0,
2153},
2154{
2155   .name = "has_input_primitive_amd",
2156   .num_srcs = 0,
2157   .has_dest = true,
2158   .dest_components = 1,
2159   .dest_bit_sizes = 0x1,
2160   .bit_size_src = -1,
2161   .num_indices = 0,
2162   .flags = 0,
2163},
2164{
2165   .name = "has_input_vertex_amd",
2166   .num_srcs = 0,
2167   .has_dest = true,
2168   .dest_components = 1,
2169   .dest_bit_sizes = 0x1,
2170   .bit_size_src = -1,
2171   .num_indices = 0,
2172   .flags = 0,
2173},
2174{
2175   .name = "ignore_ray_intersection",
2176   .num_srcs = 0,
2177   .has_dest = false,
2178   .dest_components = 0,
2179   .dest_bit_sizes = 0x0,
2180   .bit_size_src = -1,
2181   .num_indices = 0,
2182   .flags = 0,
2183},
2184{
2185   .name = "image_atomic_add",
2186   .num_srcs = 4,
2187   .src_components = {
2188      1, 4, 1, 1
2189   },
2190   .has_dest = true,
2191   .dest_components = 1,
2192   .dest_bit_sizes = 0x0,
2193   .bit_size_src = -1,
2194   .num_indices = 4,
2195   .indices = {
2196      NIR_INTRINSIC_IMAGE_DIM,
2197      NIR_INTRINSIC_IMAGE_ARRAY,
2198      NIR_INTRINSIC_FORMAT,
2199      NIR_INTRINSIC_ACCESS,
2200   },
2201   .index_map = {
2202      [NIR_INTRINSIC_IMAGE_DIM] = 1,
2203      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
2204      [NIR_INTRINSIC_FORMAT] = 3,
2205      [NIR_INTRINSIC_ACCESS] = 4,
2206    },
2207   .flags = 0,
2208},
2209{
2210   .name = "image_atomic_and",
2211   .num_srcs = 4,
2212   .src_components = {
2213      1, 4, 1, 1
2214   },
2215   .has_dest = true,
2216   .dest_components = 1,
2217   .dest_bit_sizes = 0x0,
2218   .bit_size_src = -1,
2219   .num_indices = 4,
2220   .indices = {
2221      NIR_INTRINSIC_IMAGE_DIM,
2222      NIR_INTRINSIC_IMAGE_ARRAY,
2223      NIR_INTRINSIC_FORMAT,
2224      NIR_INTRINSIC_ACCESS,
2225   },
2226   .index_map = {
2227      [NIR_INTRINSIC_IMAGE_DIM] = 1,
2228      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
2229      [NIR_INTRINSIC_FORMAT] = 3,
2230      [NIR_INTRINSIC_ACCESS] = 4,
2231    },
2232   .flags = 0,
2233},
2234{
2235   .name = "image_atomic_comp_swap",
2236   .num_srcs = 5,
2237   .src_components = {
2238      1, 4, 1, 1, 1
2239   },
2240   .has_dest = true,
2241   .dest_components = 1,
2242   .dest_bit_sizes = 0x0,
2243   .bit_size_src = -1,
2244   .num_indices = 4,
2245   .indices = {
2246      NIR_INTRINSIC_IMAGE_DIM,
2247      NIR_INTRINSIC_IMAGE_ARRAY,
2248      NIR_INTRINSIC_FORMAT,
2249      NIR_INTRINSIC_ACCESS,
2250   },
2251   .index_map = {
2252      [NIR_INTRINSIC_IMAGE_DIM] = 1,
2253      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
2254      [NIR_INTRINSIC_FORMAT] = 3,
2255      [NIR_INTRINSIC_ACCESS] = 4,
2256    },
2257   .flags = 0,
2258},
2259{
2260   .name = "image_atomic_dec_wrap",
2261   .num_srcs = 4,
2262   .src_components = {
2263      1, 4, 1, 1
2264   },
2265   .has_dest = true,
2266   .dest_components = 1,
2267   .dest_bit_sizes = 0x0,
2268   .bit_size_src = -1,
2269   .num_indices = 4,
2270   .indices = {
2271      NIR_INTRINSIC_IMAGE_DIM,
2272      NIR_INTRINSIC_IMAGE_ARRAY,
2273      NIR_INTRINSIC_FORMAT,
2274      NIR_INTRINSIC_ACCESS,
2275   },
2276   .index_map = {
2277      [NIR_INTRINSIC_IMAGE_DIM] = 1,
2278      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
2279      [NIR_INTRINSIC_FORMAT] = 3,
2280      [NIR_INTRINSIC_ACCESS] = 4,
2281    },
2282   .flags = 0,
2283},
2284{
2285   .name = "image_atomic_exchange",
2286   .num_srcs = 4,
2287   .src_components = {
2288      1, 4, 1, 1
2289   },
2290   .has_dest = true,
2291   .dest_components = 1,
2292   .dest_bit_sizes = 0x0,
2293   .bit_size_src = -1,
2294   .num_indices = 4,
2295   .indices = {
2296      NIR_INTRINSIC_IMAGE_DIM,
2297      NIR_INTRINSIC_IMAGE_ARRAY,
2298      NIR_INTRINSIC_FORMAT,
2299      NIR_INTRINSIC_ACCESS,
2300   },
2301   .index_map = {
2302      [NIR_INTRINSIC_IMAGE_DIM] = 1,
2303      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
2304      [NIR_INTRINSIC_FORMAT] = 3,
2305      [NIR_INTRINSIC_ACCESS] = 4,
2306    },
2307   .flags = 0,
2308},
2309{
2310   .name = "image_atomic_fadd",
2311   .num_srcs = 4,
2312   .src_components = {
2313      1, 4, 1, 1
2314   },
2315   .has_dest = true,
2316   .dest_components = 1,
2317   .dest_bit_sizes = 0x0,
2318   .bit_size_src = -1,
2319   .num_indices = 4,
2320   .indices = {
2321      NIR_INTRINSIC_IMAGE_DIM,
2322      NIR_INTRINSIC_IMAGE_ARRAY,
2323      NIR_INTRINSIC_FORMAT,
2324      NIR_INTRINSIC_ACCESS,
2325   },
2326   .index_map = {
2327      [NIR_INTRINSIC_IMAGE_DIM] = 1,
2328      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
2329      [NIR_INTRINSIC_FORMAT] = 3,
2330      [NIR_INTRINSIC_ACCESS] = 4,
2331    },
2332   .flags = 0,
2333},
2334{
2335   .name = "image_atomic_fmax",
2336   .num_srcs = 4,
2337   .src_components = {
2338      1, 4, 1, 1
2339   },
2340   .has_dest = true,
2341   .dest_components = 1,
2342   .dest_bit_sizes = 0x0,
2343   .bit_size_src = -1,
2344   .num_indices = 4,
2345   .indices = {
2346      NIR_INTRINSIC_IMAGE_DIM,
2347      NIR_INTRINSIC_IMAGE_ARRAY,
2348      NIR_INTRINSIC_FORMAT,
2349      NIR_INTRINSIC_ACCESS,
2350   },
2351   .index_map = {
2352      [NIR_INTRINSIC_IMAGE_DIM] = 1,
2353      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
2354      [NIR_INTRINSIC_FORMAT] = 3,
2355      [NIR_INTRINSIC_ACCESS] = 4,
2356    },
2357   .flags = 0,
2358},
2359{
2360   .name = "image_atomic_fmin",
2361   .num_srcs = 4,
2362   .src_components = {
2363      1, 4, 1, 1
2364   },
2365   .has_dest = true,
2366   .dest_components = 1,
2367   .dest_bit_sizes = 0x0,
2368   .bit_size_src = -1,
2369   .num_indices = 4,
2370   .indices = {
2371      NIR_INTRINSIC_IMAGE_DIM,
2372      NIR_INTRINSIC_IMAGE_ARRAY,
2373      NIR_INTRINSIC_FORMAT,
2374      NIR_INTRINSIC_ACCESS,
2375   },
2376   .index_map = {
2377      [NIR_INTRINSIC_IMAGE_DIM] = 1,
2378      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
2379      [NIR_INTRINSIC_FORMAT] = 3,
2380      [NIR_INTRINSIC_ACCESS] = 4,
2381    },
2382   .flags = 0,
2383},
2384{
2385   .name = "image_atomic_imax",
2386   .num_srcs = 4,
2387   .src_components = {
2388      1, 4, 1, 1
2389   },
2390   .has_dest = true,
2391   .dest_components = 1,
2392   .dest_bit_sizes = 0x0,
2393   .bit_size_src = -1,
2394   .num_indices = 4,
2395   .indices = {
2396      NIR_INTRINSIC_IMAGE_DIM,
2397      NIR_INTRINSIC_IMAGE_ARRAY,
2398      NIR_INTRINSIC_FORMAT,
2399      NIR_INTRINSIC_ACCESS,
2400   },
2401   .index_map = {
2402      [NIR_INTRINSIC_IMAGE_DIM] = 1,
2403      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
2404      [NIR_INTRINSIC_FORMAT] = 3,
2405      [NIR_INTRINSIC_ACCESS] = 4,
2406    },
2407   .flags = 0,
2408},
2409{
2410   .name = "image_atomic_imin",
2411   .num_srcs = 4,
2412   .src_components = {
2413      1, 4, 1, 1
2414   },
2415   .has_dest = true,
2416   .dest_components = 1,
2417   .dest_bit_sizes = 0x0,
2418   .bit_size_src = -1,
2419   .num_indices = 4,
2420   .indices = {
2421      NIR_INTRINSIC_IMAGE_DIM,
2422      NIR_INTRINSIC_IMAGE_ARRAY,
2423      NIR_INTRINSIC_FORMAT,
2424      NIR_INTRINSIC_ACCESS,
2425   },
2426   .index_map = {
2427      [NIR_INTRINSIC_IMAGE_DIM] = 1,
2428      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
2429      [NIR_INTRINSIC_FORMAT] = 3,
2430      [NIR_INTRINSIC_ACCESS] = 4,
2431    },
2432   .flags = 0,
2433},
2434{
2435   .name = "image_atomic_inc_wrap",
2436   .num_srcs = 4,
2437   .src_components = {
2438      1, 4, 1, 1
2439   },
2440   .has_dest = true,
2441   .dest_components = 1,
2442   .dest_bit_sizes = 0x0,
2443   .bit_size_src = -1,
2444   .num_indices = 4,
2445   .indices = {
2446      NIR_INTRINSIC_IMAGE_DIM,
2447      NIR_INTRINSIC_IMAGE_ARRAY,
2448      NIR_INTRINSIC_FORMAT,
2449      NIR_INTRINSIC_ACCESS,
2450   },
2451   .index_map = {
2452      [NIR_INTRINSIC_IMAGE_DIM] = 1,
2453      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
2454      [NIR_INTRINSIC_FORMAT] = 3,
2455      [NIR_INTRINSIC_ACCESS] = 4,
2456    },
2457   .flags = 0,
2458},
2459{
2460   .name = "image_atomic_or",
2461   .num_srcs = 4,
2462   .src_components = {
2463      1, 4, 1, 1
2464   },
2465   .has_dest = true,
2466   .dest_components = 1,
2467   .dest_bit_sizes = 0x0,
2468   .bit_size_src = -1,
2469   .num_indices = 4,
2470   .indices = {
2471      NIR_INTRINSIC_IMAGE_DIM,
2472      NIR_INTRINSIC_IMAGE_ARRAY,
2473      NIR_INTRINSIC_FORMAT,
2474      NIR_INTRINSIC_ACCESS,
2475   },
2476   .index_map = {
2477      [NIR_INTRINSIC_IMAGE_DIM] = 1,
2478      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
2479      [NIR_INTRINSIC_FORMAT] = 3,
2480      [NIR_INTRINSIC_ACCESS] = 4,
2481    },
2482   .flags = 0,
2483},
2484{
2485   .name = "image_atomic_umax",
2486   .num_srcs = 4,
2487   .src_components = {
2488      1, 4, 1, 1
2489   },
2490   .has_dest = true,
2491   .dest_components = 1,
2492   .dest_bit_sizes = 0x0,
2493   .bit_size_src = -1,
2494   .num_indices = 4,
2495   .indices = {
2496      NIR_INTRINSIC_IMAGE_DIM,
2497      NIR_INTRINSIC_IMAGE_ARRAY,
2498      NIR_INTRINSIC_FORMAT,
2499      NIR_INTRINSIC_ACCESS,
2500   },
2501   .index_map = {
2502      [NIR_INTRINSIC_IMAGE_DIM] = 1,
2503      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
2504      [NIR_INTRINSIC_FORMAT] = 3,
2505      [NIR_INTRINSIC_ACCESS] = 4,
2506    },
2507   .flags = 0,
2508},
2509{
2510   .name = "image_atomic_umin",
2511   .num_srcs = 4,
2512   .src_components = {
2513      1, 4, 1, 1
2514   },
2515   .has_dest = true,
2516   .dest_components = 1,
2517   .dest_bit_sizes = 0x0,
2518   .bit_size_src = -1,
2519   .num_indices = 4,
2520   .indices = {
2521      NIR_INTRINSIC_IMAGE_DIM,
2522      NIR_INTRINSIC_IMAGE_ARRAY,
2523      NIR_INTRINSIC_FORMAT,
2524      NIR_INTRINSIC_ACCESS,
2525   },
2526   .index_map = {
2527      [NIR_INTRINSIC_IMAGE_DIM] = 1,
2528      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
2529      [NIR_INTRINSIC_FORMAT] = 3,
2530      [NIR_INTRINSIC_ACCESS] = 4,
2531    },
2532   .flags = 0,
2533},
2534{
2535   .name = "image_atomic_xor",
2536   .num_srcs = 4,
2537   .src_components = {
2538      1, 4, 1, 1
2539   },
2540   .has_dest = true,
2541   .dest_components = 1,
2542   .dest_bit_sizes = 0x0,
2543   .bit_size_src = -1,
2544   .num_indices = 4,
2545   .indices = {
2546      NIR_INTRINSIC_IMAGE_DIM,
2547      NIR_INTRINSIC_IMAGE_ARRAY,
2548      NIR_INTRINSIC_FORMAT,
2549      NIR_INTRINSIC_ACCESS,
2550   },
2551   .index_map = {
2552      [NIR_INTRINSIC_IMAGE_DIM] = 1,
2553      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
2554      [NIR_INTRINSIC_FORMAT] = 3,
2555      [NIR_INTRINSIC_ACCESS] = 4,
2556    },
2557   .flags = 0,
2558},
2559{
2560   .name = "image_deref_atomic_add",
2561   .num_srcs = 4,
2562   .src_components = {
2563      -1, 4, 1, 1
2564   },
2565   .has_dest = true,
2566   .dest_components = 1,
2567   .dest_bit_sizes = 0x0,
2568   .bit_size_src = -1,
2569   .num_indices = 4,
2570   .indices = {
2571      NIR_INTRINSIC_IMAGE_DIM,
2572      NIR_INTRINSIC_IMAGE_ARRAY,
2573      NIR_INTRINSIC_FORMAT,
2574      NIR_INTRINSIC_ACCESS,
2575   },
2576   .index_map = {
2577      [NIR_INTRINSIC_IMAGE_DIM] = 1,
2578      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
2579      [NIR_INTRINSIC_FORMAT] = 3,
2580      [NIR_INTRINSIC_ACCESS] = 4,
2581    },
2582   .flags = 0,
2583},
2584{
2585   .name = "image_deref_atomic_and",
2586   .num_srcs = 4,
2587   .src_components = {
2588      -1, 4, 1, 1
2589   },
2590   .has_dest = true,
2591   .dest_components = 1,
2592   .dest_bit_sizes = 0x0,
2593   .bit_size_src = -1,
2594   .num_indices = 4,
2595   .indices = {
2596      NIR_INTRINSIC_IMAGE_DIM,
2597      NIR_INTRINSIC_IMAGE_ARRAY,
2598      NIR_INTRINSIC_FORMAT,
2599      NIR_INTRINSIC_ACCESS,
2600   },
2601   .index_map = {
2602      [NIR_INTRINSIC_IMAGE_DIM] = 1,
2603      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
2604      [NIR_INTRINSIC_FORMAT] = 3,
2605      [NIR_INTRINSIC_ACCESS] = 4,
2606    },
2607   .flags = 0,
2608},
2609{
2610   .name = "image_deref_atomic_comp_swap",
2611   .num_srcs = 5,
2612   .src_components = {
2613      -1, 4, 1, 1, 1
2614   },
2615   .has_dest = true,
2616   .dest_components = 1,
2617   .dest_bit_sizes = 0x0,
2618   .bit_size_src = -1,
2619   .num_indices = 4,
2620   .indices = {
2621      NIR_INTRINSIC_IMAGE_DIM,
2622      NIR_INTRINSIC_IMAGE_ARRAY,
2623      NIR_INTRINSIC_FORMAT,
2624      NIR_INTRINSIC_ACCESS,
2625   },
2626   .index_map = {
2627      [NIR_INTRINSIC_IMAGE_DIM] = 1,
2628      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
2629      [NIR_INTRINSIC_FORMAT] = 3,
2630      [NIR_INTRINSIC_ACCESS] = 4,
2631    },
2632   .flags = 0,
2633},
2634{
2635   .name = "image_deref_atomic_dec_wrap",
2636   .num_srcs = 4,
2637   .src_components = {
2638      -1, 4, 1, 1
2639   },
2640   .has_dest = true,
2641   .dest_components = 1,
2642   .dest_bit_sizes = 0x0,
2643   .bit_size_src = -1,
2644   .num_indices = 4,
2645   .indices = {
2646      NIR_INTRINSIC_IMAGE_DIM,
2647      NIR_INTRINSIC_IMAGE_ARRAY,
2648      NIR_INTRINSIC_FORMAT,
2649      NIR_INTRINSIC_ACCESS,
2650   },
2651   .index_map = {
2652      [NIR_INTRINSIC_IMAGE_DIM] = 1,
2653      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
2654      [NIR_INTRINSIC_FORMAT] = 3,
2655      [NIR_INTRINSIC_ACCESS] = 4,
2656    },
2657   .flags = 0,
2658},
2659{
2660   .name = "image_deref_atomic_exchange",
2661   .num_srcs = 4,
2662   .src_components = {
2663      -1, 4, 1, 1
2664   },
2665   .has_dest = true,
2666   .dest_components = 1,
2667   .dest_bit_sizes = 0x0,
2668   .bit_size_src = -1,
2669   .num_indices = 4,
2670   .indices = {
2671      NIR_INTRINSIC_IMAGE_DIM,
2672      NIR_INTRINSIC_IMAGE_ARRAY,
2673      NIR_INTRINSIC_FORMAT,
2674      NIR_INTRINSIC_ACCESS,
2675   },
2676   .index_map = {
2677      [NIR_INTRINSIC_IMAGE_DIM] = 1,
2678      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
2679      [NIR_INTRINSIC_FORMAT] = 3,
2680      [NIR_INTRINSIC_ACCESS] = 4,
2681    },
2682   .flags = 0,
2683},
2684{
2685   .name = "image_deref_atomic_fadd",
2686   .num_srcs = 4,
2687   .src_components = {
2688      -1, 4, 1, 1
2689   },
2690   .has_dest = true,
2691   .dest_components = 1,
2692   .dest_bit_sizes = 0x0,
2693   .bit_size_src = -1,
2694   .num_indices = 4,
2695   .indices = {
2696      NIR_INTRINSIC_IMAGE_DIM,
2697      NIR_INTRINSIC_IMAGE_ARRAY,
2698      NIR_INTRINSIC_FORMAT,
2699      NIR_INTRINSIC_ACCESS,
2700   },
2701   .index_map = {
2702      [NIR_INTRINSIC_IMAGE_DIM] = 1,
2703      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
2704      [NIR_INTRINSIC_FORMAT] = 3,
2705      [NIR_INTRINSIC_ACCESS] = 4,
2706    },
2707   .flags = 0,
2708},
2709{
2710   .name = "image_deref_atomic_fmax",
2711   .num_srcs = 4,
2712   .src_components = {
2713      -1, 4, 1, 1
2714   },
2715   .has_dest = true,
2716   .dest_components = 1,
2717   .dest_bit_sizes = 0x0,
2718   .bit_size_src = -1,
2719   .num_indices = 4,
2720   .indices = {
2721      NIR_INTRINSIC_IMAGE_DIM,
2722      NIR_INTRINSIC_IMAGE_ARRAY,
2723      NIR_INTRINSIC_FORMAT,
2724      NIR_INTRINSIC_ACCESS,
2725   },
2726   .index_map = {
2727      [NIR_INTRINSIC_IMAGE_DIM] = 1,
2728      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
2729      [NIR_INTRINSIC_FORMAT] = 3,
2730      [NIR_INTRINSIC_ACCESS] = 4,
2731    },
2732   .flags = 0,
2733},
2734{
2735   .name = "image_deref_atomic_fmin",
2736   .num_srcs = 4,
2737   .src_components = {
2738      -1, 4, 1, 1
2739   },
2740   .has_dest = true,
2741   .dest_components = 1,
2742   .dest_bit_sizes = 0x0,
2743   .bit_size_src = -1,
2744   .num_indices = 4,
2745   .indices = {
2746      NIR_INTRINSIC_IMAGE_DIM,
2747      NIR_INTRINSIC_IMAGE_ARRAY,
2748      NIR_INTRINSIC_FORMAT,
2749      NIR_INTRINSIC_ACCESS,
2750   },
2751   .index_map = {
2752      [NIR_INTRINSIC_IMAGE_DIM] = 1,
2753      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
2754      [NIR_INTRINSIC_FORMAT] = 3,
2755      [NIR_INTRINSIC_ACCESS] = 4,
2756    },
2757   .flags = 0,
2758},
2759{
2760   .name = "image_deref_atomic_imax",
2761   .num_srcs = 4,
2762   .src_components = {
2763      -1, 4, 1, 1
2764   },
2765   .has_dest = true,
2766   .dest_components = 1,
2767   .dest_bit_sizes = 0x0,
2768   .bit_size_src = -1,
2769   .num_indices = 4,
2770   .indices = {
2771      NIR_INTRINSIC_IMAGE_DIM,
2772      NIR_INTRINSIC_IMAGE_ARRAY,
2773      NIR_INTRINSIC_FORMAT,
2774      NIR_INTRINSIC_ACCESS,
2775   },
2776   .index_map = {
2777      [NIR_INTRINSIC_IMAGE_DIM] = 1,
2778      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
2779      [NIR_INTRINSIC_FORMAT] = 3,
2780      [NIR_INTRINSIC_ACCESS] = 4,
2781    },
2782   .flags = 0,
2783},
2784{
2785   .name = "image_deref_atomic_imin",
2786   .num_srcs = 4,
2787   .src_components = {
2788      -1, 4, 1, 1
2789   },
2790   .has_dest = true,
2791   .dest_components = 1,
2792   .dest_bit_sizes = 0x0,
2793   .bit_size_src = -1,
2794   .num_indices = 4,
2795   .indices = {
2796      NIR_INTRINSIC_IMAGE_DIM,
2797      NIR_INTRINSIC_IMAGE_ARRAY,
2798      NIR_INTRINSIC_FORMAT,
2799      NIR_INTRINSIC_ACCESS,
2800   },
2801   .index_map = {
2802      [NIR_INTRINSIC_IMAGE_DIM] = 1,
2803      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
2804      [NIR_INTRINSIC_FORMAT] = 3,
2805      [NIR_INTRINSIC_ACCESS] = 4,
2806    },
2807   .flags = 0,
2808},
2809{
2810   .name = "image_deref_atomic_inc_wrap",
2811   .num_srcs = 4,
2812   .src_components = {
2813      -1, 4, 1, 1
2814   },
2815   .has_dest = true,
2816   .dest_components = 1,
2817   .dest_bit_sizes = 0x0,
2818   .bit_size_src = -1,
2819   .num_indices = 4,
2820   .indices = {
2821      NIR_INTRINSIC_IMAGE_DIM,
2822      NIR_INTRINSIC_IMAGE_ARRAY,
2823      NIR_INTRINSIC_FORMAT,
2824      NIR_INTRINSIC_ACCESS,
2825   },
2826   .index_map = {
2827      [NIR_INTRINSIC_IMAGE_DIM] = 1,
2828      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
2829      [NIR_INTRINSIC_FORMAT] = 3,
2830      [NIR_INTRINSIC_ACCESS] = 4,
2831    },
2832   .flags = 0,
2833},
2834{
2835   .name = "image_deref_atomic_or",
2836   .num_srcs = 4,
2837   .src_components = {
2838      -1, 4, 1, 1
2839   },
2840   .has_dest = true,
2841   .dest_components = 1,
2842   .dest_bit_sizes = 0x0,
2843   .bit_size_src = -1,
2844   .num_indices = 4,
2845   .indices = {
2846      NIR_INTRINSIC_IMAGE_DIM,
2847      NIR_INTRINSIC_IMAGE_ARRAY,
2848      NIR_INTRINSIC_FORMAT,
2849      NIR_INTRINSIC_ACCESS,
2850   },
2851   .index_map = {
2852      [NIR_INTRINSIC_IMAGE_DIM] = 1,
2853      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
2854      [NIR_INTRINSIC_FORMAT] = 3,
2855      [NIR_INTRINSIC_ACCESS] = 4,
2856    },
2857   .flags = 0,
2858},
2859{
2860   .name = "image_deref_atomic_umax",
2861   .num_srcs = 4,
2862   .src_components = {
2863      -1, 4, 1, 1
2864   },
2865   .has_dest = true,
2866   .dest_components = 1,
2867   .dest_bit_sizes = 0x0,
2868   .bit_size_src = -1,
2869   .num_indices = 4,
2870   .indices = {
2871      NIR_INTRINSIC_IMAGE_DIM,
2872      NIR_INTRINSIC_IMAGE_ARRAY,
2873      NIR_INTRINSIC_FORMAT,
2874      NIR_INTRINSIC_ACCESS,
2875   },
2876   .index_map = {
2877      [NIR_INTRINSIC_IMAGE_DIM] = 1,
2878      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
2879      [NIR_INTRINSIC_FORMAT] = 3,
2880      [NIR_INTRINSIC_ACCESS] = 4,
2881    },
2882   .flags = 0,
2883},
2884{
2885   .name = "image_deref_atomic_umin",
2886   .num_srcs = 4,
2887   .src_components = {
2888      -1, 4, 1, 1
2889   },
2890   .has_dest = true,
2891   .dest_components = 1,
2892   .dest_bit_sizes = 0x0,
2893   .bit_size_src = -1,
2894   .num_indices = 4,
2895   .indices = {
2896      NIR_INTRINSIC_IMAGE_DIM,
2897      NIR_INTRINSIC_IMAGE_ARRAY,
2898      NIR_INTRINSIC_FORMAT,
2899      NIR_INTRINSIC_ACCESS,
2900   },
2901   .index_map = {
2902      [NIR_INTRINSIC_IMAGE_DIM] = 1,
2903      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
2904      [NIR_INTRINSIC_FORMAT] = 3,
2905      [NIR_INTRINSIC_ACCESS] = 4,
2906    },
2907   .flags = 0,
2908},
2909{
2910   .name = "image_deref_atomic_xor",
2911   .num_srcs = 4,
2912   .src_components = {
2913      -1, 4, 1, 1
2914   },
2915   .has_dest = true,
2916   .dest_components = 1,
2917   .dest_bit_sizes = 0x0,
2918   .bit_size_src = -1,
2919   .num_indices = 4,
2920   .indices = {
2921      NIR_INTRINSIC_IMAGE_DIM,
2922      NIR_INTRINSIC_IMAGE_ARRAY,
2923      NIR_INTRINSIC_FORMAT,
2924      NIR_INTRINSIC_ACCESS,
2925   },
2926   .index_map = {
2927      [NIR_INTRINSIC_IMAGE_DIM] = 1,
2928      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
2929      [NIR_INTRINSIC_FORMAT] = 3,
2930      [NIR_INTRINSIC_ACCESS] = 4,
2931    },
2932   .flags = 0,
2933},
2934{
2935   .name = "image_deref_format",
2936   .num_srcs = 1,
2937   .src_components = {
2938      -1
2939   },
2940   .has_dest = true,
2941   .dest_components = 1,
2942   .dest_bit_sizes = 0x0,
2943   .bit_size_src = -1,
2944   .num_indices = 4,
2945   .indices = {
2946      NIR_INTRINSIC_IMAGE_DIM,
2947      NIR_INTRINSIC_IMAGE_ARRAY,
2948      NIR_INTRINSIC_FORMAT,
2949      NIR_INTRINSIC_ACCESS,
2950   },
2951   .index_map = {
2952      [NIR_INTRINSIC_IMAGE_DIM] = 1,
2953      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
2954      [NIR_INTRINSIC_FORMAT] = 3,
2955      [NIR_INTRINSIC_ACCESS] = 4,
2956    },
2957   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
2958},
2959{
2960   .name = "image_deref_load",
2961   .num_srcs = 4,
2962   .src_components = {
2963      -1, 4, 1, 1
2964   },
2965   .has_dest = true,
2966   .dest_components = 0,
2967   .dest_bit_sizes = 0x0,
2968   .bit_size_src = -1,
2969   .num_indices = 5,
2970   .indices = {
2971      NIR_INTRINSIC_IMAGE_DIM,
2972      NIR_INTRINSIC_IMAGE_ARRAY,
2973      NIR_INTRINSIC_FORMAT,
2974      NIR_INTRINSIC_ACCESS,
2975      NIR_INTRINSIC_DEST_TYPE,
2976   },
2977   .index_map = {
2978      [NIR_INTRINSIC_IMAGE_DIM] = 1,
2979      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
2980      [NIR_INTRINSIC_FORMAT] = 3,
2981      [NIR_INTRINSIC_ACCESS] = 4,
2982      [NIR_INTRINSIC_DEST_TYPE] = 5,
2983    },
2984   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
2985},
2986{
2987   .name = "image_deref_load_param_intel",
2988   .num_srcs = 1,
2989   .src_components = {
2990      1
2991   },
2992   .has_dest = true,
2993   .dest_components = 0,
2994   .dest_bit_sizes = 0x0,
2995   .bit_size_src = -1,
2996   .num_indices = 1,
2997   .indices = {
2998      NIR_INTRINSIC_BASE,
2999   },
3000   .index_map = {
3001      [NIR_INTRINSIC_BASE] = 1,
3002    },
3003   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3004},
3005{
3006   .name = "image_deref_load_raw_intel",
3007   .num_srcs = 2,
3008   .src_components = {
3009      -1, 1
3010   },
3011   .has_dest = true,
3012   .dest_components = 0,
3013   .dest_bit_sizes = 0x0,
3014   .bit_size_src = -1,
3015   .num_indices = 4,
3016   .indices = {
3017      NIR_INTRINSIC_IMAGE_DIM,
3018      NIR_INTRINSIC_IMAGE_ARRAY,
3019      NIR_INTRINSIC_FORMAT,
3020      NIR_INTRINSIC_ACCESS,
3021   },
3022   .index_map = {
3023      [NIR_INTRINSIC_IMAGE_DIM] = 1,
3024      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
3025      [NIR_INTRINSIC_FORMAT] = 3,
3026      [NIR_INTRINSIC_ACCESS] = 4,
3027    },
3028   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
3029},
3030{
3031   .name = "image_deref_order",
3032   .num_srcs = 1,
3033   .src_components = {
3034      -1
3035   },
3036   .has_dest = true,
3037   .dest_components = 1,
3038   .dest_bit_sizes = 0x0,
3039   .bit_size_src = -1,
3040   .num_indices = 4,
3041   .indices = {
3042      NIR_INTRINSIC_IMAGE_DIM,
3043      NIR_INTRINSIC_IMAGE_ARRAY,
3044      NIR_INTRINSIC_FORMAT,
3045      NIR_INTRINSIC_ACCESS,
3046   },
3047   .index_map = {
3048      [NIR_INTRINSIC_IMAGE_DIM] = 1,
3049      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
3050      [NIR_INTRINSIC_FORMAT] = 3,
3051      [NIR_INTRINSIC_ACCESS] = 4,
3052    },
3053   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3054},
3055{
3056   .name = "image_deref_samples",
3057   .num_srcs = 1,
3058   .src_components = {
3059      -1
3060   },
3061   .has_dest = true,
3062   .dest_components = 1,
3063   .dest_bit_sizes = 0x0,
3064   .bit_size_src = -1,
3065   .num_indices = 4,
3066   .indices = {
3067      NIR_INTRINSIC_IMAGE_DIM,
3068      NIR_INTRINSIC_IMAGE_ARRAY,
3069      NIR_INTRINSIC_FORMAT,
3070      NIR_INTRINSIC_ACCESS,
3071   },
3072   .index_map = {
3073      [NIR_INTRINSIC_IMAGE_DIM] = 1,
3074      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
3075      [NIR_INTRINSIC_FORMAT] = 3,
3076      [NIR_INTRINSIC_ACCESS] = 4,
3077    },
3078   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3079},
3080{
3081   .name = "image_deref_size",
3082   .num_srcs = 2,
3083   .src_components = {
3084      -1, 1
3085   },
3086   .has_dest = true,
3087   .dest_components = 0,
3088   .dest_bit_sizes = 0x0,
3089   .bit_size_src = -1,
3090   .num_indices = 4,
3091   .indices = {
3092      NIR_INTRINSIC_IMAGE_DIM,
3093      NIR_INTRINSIC_IMAGE_ARRAY,
3094      NIR_INTRINSIC_FORMAT,
3095      NIR_INTRINSIC_ACCESS,
3096   },
3097   .index_map = {
3098      [NIR_INTRINSIC_IMAGE_DIM] = 1,
3099      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
3100      [NIR_INTRINSIC_FORMAT] = 3,
3101      [NIR_INTRINSIC_ACCESS] = 4,
3102    },
3103   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3104},
3105{
3106   .name = "image_deref_sparse_load",
3107   .num_srcs = 4,
3108   .src_components = {
3109      -1, 4, 1, 1
3110   },
3111   .has_dest = true,
3112   .dest_components = 0,
3113   .dest_bit_sizes = 0x0,
3114   .bit_size_src = -1,
3115   .num_indices = 5,
3116   .indices = {
3117      NIR_INTRINSIC_IMAGE_DIM,
3118      NIR_INTRINSIC_IMAGE_ARRAY,
3119      NIR_INTRINSIC_FORMAT,
3120      NIR_INTRINSIC_ACCESS,
3121      NIR_INTRINSIC_DEST_TYPE,
3122   },
3123   .index_map = {
3124      [NIR_INTRINSIC_IMAGE_DIM] = 1,
3125      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
3126      [NIR_INTRINSIC_FORMAT] = 3,
3127      [NIR_INTRINSIC_ACCESS] = 4,
3128      [NIR_INTRINSIC_DEST_TYPE] = 5,
3129    },
3130   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
3131},
3132{
3133   .name = "image_deref_store",
3134   .num_srcs = 5,
3135   .src_components = {
3136      -1, 4, 1, 0, 1
3137   },
3138   .has_dest = false,
3139   .dest_components = 0,
3140   .dest_bit_sizes = 0x0,
3141   .bit_size_src = -1,
3142   .num_indices = 5,
3143   .indices = {
3144      NIR_INTRINSIC_IMAGE_DIM,
3145      NIR_INTRINSIC_IMAGE_ARRAY,
3146      NIR_INTRINSIC_FORMAT,
3147      NIR_INTRINSIC_ACCESS,
3148      NIR_INTRINSIC_SRC_TYPE,
3149   },
3150   .index_map = {
3151      [NIR_INTRINSIC_IMAGE_DIM] = 1,
3152      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
3153      [NIR_INTRINSIC_FORMAT] = 3,
3154      [NIR_INTRINSIC_ACCESS] = 4,
3155      [NIR_INTRINSIC_SRC_TYPE] = 5,
3156    },
3157   .flags = 0,
3158},
3159{
3160   .name = "image_deref_store_raw_intel",
3161   .num_srcs = 3,
3162   .src_components = {
3163      -1, 1, 0
3164   },
3165   .has_dest = false,
3166   .dest_components = 0,
3167   .dest_bit_sizes = 0x0,
3168   .bit_size_src = -1,
3169   .num_indices = 4,
3170   .indices = {
3171      NIR_INTRINSIC_IMAGE_DIM,
3172      NIR_INTRINSIC_IMAGE_ARRAY,
3173      NIR_INTRINSIC_FORMAT,
3174      NIR_INTRINSIC_ACCESS,
3175   },
3176   .index_map = {
3177      [NIR_INTRINSIC_IMAGE_DIM] = 1,
3178      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
3179      [NIR_INTRINSIC_FORMAT] = 3,
3180      [NIR_INTRINSIC_ACCESS] = 4,
3181    },
3182   .flags = 0,
3183},
3184{
3185   .name = "image_format",
3186   .num_srcs = 1,
3187   .src_components = {
3188      1
3189   },
3190   .has_dest = true,
3191   .dest_components = 1,
3192   .dest_bit_sizes = 0x0,
3193   .bit_size_src = -1,
3194   .num_indices = 4,
3195   .indices = {
3196      NIR_INTRINSIC_IMAGE_DIM,
3197      NIR_INTRINSIC_IMAGE_ARRAY,
3198      NIR_INTRINSIC_FORMAT,
3199      NIR_INTRINSIC_ACCESS,
3200   },
3201   .index_map = {
3202      [NIR_INTRINSIC_IMAGE_DIM] = 1,
3203      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
3204      [NIR_INTRINSIC_FORMAT] = 3,
3205      [NIR_INTRINSIC_ACCESS] = 4,
3206    },
3207   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3208},
3209{
3210   .name = "image_load",
3211   .num_srcs = 4,
3212   .src_components = {
3213      1, 4, 1, 1
3214   },
3215   .has_dest = true,
3216   .dest_components = 0,
3217   .dest_bit_sizes = 0x0,
3218   .bit_size_src = -1,
3219   .num_indices = 5,
3220   .indices = {
3221      NIR_INTRINSIC_IMAGE_DIM,
3222      NIR_INTRINSIC_IMAGE_ARRAY,
3223      NIR_INTRINSIC_FORMAT,
3224      NIR_INTRINSIC_ACCESS,
3225      NIR_INTRINSIC_DEST_TYPE,
3226   },
3227   .index_map = {
3228      [NIR_INTRINSIC_IMAGE_DIM] = 1,
3229      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
3230      [NIR_INTRINSIC_FORMAT] = 3,
3231      [NIR_INTRINSIC_ACCESS] = 4,
3232      [NIR_INTRINSIC_DEST_TYPE] = 5,
3233    },
3234   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
3235},
3236{
3237   .name = "image_load_raw_intel",
3238   .num_srcs = 2,
3239   .src_components = {
3240      1, 1
3241   },
3242   .has_dest = true,
3243   .dest_components = 0,
3244   .dest_bit_sizes = 0x0,
3245   .bit_size_src = -1,
3246   .num_indices = 4,
3247   .indices = {
3248      NIR_INTRINSIC_IMAGE_DIM,
3249      NIR_INTRINSIC_IMAGE_ARRAY,
3250      NIR_INTRINSIC_FORMAT,
3251      NIR_INTRINSIC_ACCESS,
3252   },
3253   .index_map = {
3254      [NIR_INTRINSIC_IMAGE_DIM] = 1,
3255      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
3256      [NIR_INTRINSIC_FORMAT] = 3,
3257      [NIR_INTRINSIC_ACCESS] = 4,
3258    },
3259   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
3260},
3261{
3262   .name = "image_order",
3263   .num_srcs = 1,
3264   .src_components = {
3265      1
3266   },
3267   .has_dest = true,
3268   .dest_components = 1,
3269   .dest_bit_sizes = 0x0,
3270   .bit_size_src = -1,
3271   .num_indices = 4,
3272   .indices = {
3273      NIR_INTRINSIC_IMAGE_DIM,
3274      NIR_INTRINSIC_IMAGE_ARRAY,
3275      NIR_INTRINSIC_FORMAT,
3276      NIR_INTRINSIC_ACCESS,
3277   },
3278   .index_map = {
3279      [NIR_INTRINSIC_IMAGE_DIM] = 1,
3280      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
3281      [NIR_INTRINSIC_FORMAT] = 3,
3282      [NIR_INTRINSIC_ACCESS] = 4,
3283    },
3284   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3285},
3286{
3287   .name = "image_samples",
3288   .num_srcs = 1,
3289   .src_components = {
3290      1
3291   },
3292   .has_dest = true,
3293   .dest_components = 1,
3294   .dest_bit_sizes = 0x0,
3295   .bit_size_src = -1,
3296   .num_indices = 4,
3297   .indices = {
3298      NIR_INTRINSIC_IMAGE_DIM,
3299      NIR_INTRINSIC_IMAGE_ARRAY,
3300      NIR_INTRINSIC_FORMAT,
3301      NIR_INTRINSIC_ACCESS,
3302   },
3303   .index_map = {
3304      [NIR_INTRINSIC_IMAGE_DIM] = 1,
3305      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
3306      [NIR_INTRINSIC_FORMAT] = 3,
3307      [NIR_INTRINSIC_ACCESS] = 4,
3308    },
3309   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3310},
3311{
3312   .name = "image_size",
3313   .num_srcs = 2,
3314   .src_components = {
3315      1, 1
3316   },
3317   .has_dest = true,
3318   .dest_components = 0,
3319   .dest_bit_sizes = 0x0,
3320   .bit_size_src = -1,
3321   .num_indices = 4,
3322   .indices = {
3323      NIR_INTRINSIC_IMAGE_DIM,
3324      NIR_INTRINSIC_IMAGE_ARRAY,
3325      NIR_INTRINSIC_FORMAT,
3326      NIR_INTRINSIC_ACCESS,
3327   },
3328   .index_map = {
3329      [NIR_INTRINSIC_IMAGE_DIM] = 1,
3330      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
3331      [NIR_INTRINSIC_FORMAT] = 3,
3332      [NIR_INTRINSIC_ACCESS] = 4,
3333    },
3334   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3335},
3336{
3337   .name = "image_sparse_load",
3338   .num_srcs = 4,
3339   .src_components = {
3340      1, 4, 1, 1
3341   },
3342   .has_dest = true,
3343   .dest_components = 0,
3344   .dest_bit_sizes = 0x0,
3345   .bit_size_src = -1,
3346   .num_indices = 5,
3347   .indices = {
3348      NIR_INTRINSIC_IMAGE_DIM,
3349      NIR_INTRINSIC_IMAGE_ARRAY,
3350      NIR_INTRINSIC_FORMAT,
3351      NIR_INTRINSIC_ACCESS,
3352      NIR_INTRINSIC_DEST_TYPE,
3353   },
3354   .index_map = {
3355      [NIR_INTRINSIC_IMAGE_DIM] = 1,
3356      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
3357      [NIR_INTRINSIC_FORMAT] = 3,
3358      [NIR_INTRINSIC_ACCESS] = 4,
3359      [NIR_INTRINSIC_DEST_TYPE] = 5,
3360    },
3361   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
3362},
3363{
3364   .name = "image_store",
3365   .num_srcs = 5,
3366   .src_components = {
3367      1, 4, 1, 0, 1
3368   },
3369   .has_dest = false,
3370   .dest_components = 0,
3371   .dest_bit_sizes = 0x0,
3372   .bit_size_src = -1,
3373   .num_indices = 5,
3374   .indices = {
3375      NIR_INTRINSIC_IMAGE_DIM,
3376      NIR_INTRINSIC_IMAGE_ARRAY,
3377      NIR_INTRINSIC_FORMAT,
3378      NIR_INTRINSIC_ACCESS,
3379      NIR_INTRINSIC_SRC_TYPE,
3380   },
3381   .index_map = {
3382      [NIR_INTRINSIC_IMAGE_DIM] = 1,
3383      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
3384      [NIR_INTRINSIC_FORMAT] = 3,
3385      [NIR_INTRINSIC_ACCESS] = 4,
3386      [NIR_INTRINSIC_SRC_TYPE] = 5,
3387    },
3388   .flags = 0,
3389},
3390{
3391   .name = "image_store_raw_intel",
3392   .num_srcs = 3,
3393   .src_components = {
3394      1, 1, 0
3395   },
3396   .has_dest = false,
3397   .dest_components = 0,
3398   .dest_bit_sizes = 0x0,
3399   .bit_size_src = -1,
3400   .num_indices = 4,
3401   .indices = {
3402      NIR_INTRINSIC_IMAGE_DIM,
3403      NIR_INTRINSIC_IMAGE_ARRAY,
3404      NIR_INTRINSIC_FORMAT,
3405      NIR_INTRINSIC_ACCESS,
3406   },
3407   .index_map = {
3408      [NIR_INTRINSIC_IMAGE_DIM] = 1,
3409      [NIR_INTRINSIC_IMAGE_ARRAY] = 2,
3410      [NIR_INTRINSIC_FORMAT] = 3,
3411      [NIR_INTRINSIC_ACCESS] = 4,
3412    },
3413   .flags = 0,
3414},
3415{
3416   .name = "inclusive_scan",
3417   .num_srcs = 1,
3418   .src_components = {
3419      0
3420   },
3421   .has_dest = true,
3422   .dest_components = 0,
3423   .dest_bit_sizes = 0x0,
3424   .bit_size_src = 0,
3425   .num_indices = 1,
3426   .indices = {
3427      NIR_INTRINSIC_REDUCTION_OP,
3428   },
3429   .index_map = {
3430      [NIR_INTRINSIC_REDUCTION_OP] = 1,
3431    },
3432   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
3433},
3434{
3435   .name = "interp_deref_at_centroid",
3436   .num_srcs = 1,
3437   .src_components = {
3438      1
3439   },
3440   .has_dest = true,
3441   .dest_components = 0,
3442   .dest_bit_sizes = 0x0,
3443   .bit_size_src = -1,
3444   .num_indices = 0,
3445   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3446},
3447{
3448   .name = "interp_deref_at_offset",
3449   .num_srcs = 2,
3450   .src_components = {
3451      1, 2
3452   },
3453   .has_dest = true,
3454   .dest_components = 0,
3455   .dest_bit_sizes = 0x0,
3456   .bit_size_src = -1,
3457   .num_indices = 0,
3458   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3459},
3460{
3461   .name = "interp_deref_at_sample",
3462   .num_srcs = 2,
3463   .src_components = {
3464      1, 1
3465   },
3466   .has_dest = true,
3467   .dest_components = 0,
3468   .dest_bit_sizes = 0x0,
3469   .bit_size_src = -1,
3470   .num_indices = 0,
3471   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3472},
3473{
3474   .name = "interp_deref_at_vertex",
3475   .num_srcs = 2,
3476   .src_components = {
3477      1, 1
3478   },
3479   .has_dest = true,
3480   .dest_components = 0,
3481   .dest_bit_sizes = 0x0,
3482   .bit_size_src = -1,
3483   .num_indices = 0,
3484   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3485},
3486{
3487   .name = "is_helper_invocation",
3488   .num_srcs = 0,
3489   .has_dest = true,
3490   .dest_components = 1,
3491   .dest_bit_sizes = 0x0,
3492   .bit_size_src = -1,
3493   .num_indices = 0,
3494   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
3495},
3496{
3497   .name = "is_sparse_texels_resident",
3498   .num_srcs = 1,
3499   .src_components = {
3500      1
3501   },
3502   .has_dest = true,
3503   .dest_components = 1,
3504   .dest_bit_sizes = 0x1,
3505   .bit_size_src = -1,
3506   .num_indices = 0,
3507   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3508},
3509{
3510   .name = "lane_permute_16_amd",
3511   .num_srcs = 3,
3512   .src_components = {
3513      1, 1, 1
3514   },
3515   .has_dest = true,
3516   .dest_components = 1,
3517   .dest_bit_sizes = 0x20,
3518   .bit_size_src = -1,
3519   .num_indices = 0,
3520   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
3521},
3522{
3523   .name = "last_invocation",
3524   .num_srcs = 0,
3525   .has_dest = true,
3526   .dest_components = 1,
3527   .dest_bit_sizes = 0x20,
3528   .bit_size_src = -1,
3529   .num_indices = 0,
3530   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
3531},
3532{
3533   .name = "load_aa_line_width",
3534   .num_srcs = 0,
3535   .has_dest = true,
3536   .dest_components = 1,
3537   .dest_bit_sizes = 0x20,
3538   .bit_size_src = -1,
3539   .num_indices = 0,
3540   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3541},
3542{
3543   .name = "load_back_face_agx",
3544   .num_srcs = 0,
3545   .has_dest = true,
3546   .dest_components = 1,
3547   .dest_bit_sizes = 0x21,
3548   .bit_size_src = -1,
3549   .num_indices = 0,
3550   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3551},
3552{
3553   .name = "load_barycentric_at_offset",
3554   .num_srcs = 1,
3555   .src_components = {
3556      2
3557   },
3558   .has_dest = true,
3559   .dest_components = 2,
3560   .dest_bit_sizes = 0x0,
3561   .bit_size_src = -1,
3562   .num_indices = 1,
3563   .indices = {
3564      NIR_INTRINSIC_INTERP_MODE,
3565   },
3566   .index_map = {
3567      [NIR_INTRINSIC_INTERP_MODE] = 1,
3568    },
3569   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3570},
3571{
3572   .name = "load_barycentric_at_sample",
3573   .num_srcs = 1,
3574   .src_components = {
3575      1
3576   },
3577   .has_dest = true,
3578   .dest_components = 2,
3579   .dest_bit_sizes = 0x0,
3580   .bit_size_src = -1,
3581   .num_indices = 1,
3582   .indices = {
3583      NIR_INTRINSIC_INTERP_MODE,
3584   },
3585   .index_map = {
3586      [NIR_INTRINSIC_INTERP_MODE] = 1,
3587    },
3588   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3589},
3590{
3591   .name = "load_barycentric_centroid",
3592   .num_srcs = 0,
3593   .has_dest = true,
3594   .dest_components = 2,
3595   .dest_bit_sizes = 0x0,
3596   .bit_size_src = -1,
3597   .num_indices = 1,
3598   .indices = {
3599      NIR_INTRINSIC_INTERP_MODE,
3600   },
3601   .index_map = {
3602      [NIR_INTRINSIC_INTERP_MODE] = 1,
3603    },
3604   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3605},
3606{
3607   .name = "load_barycentric_model",
3608   .num_srcs = 0,
3609   .has_dest = true,
3610   .dest_components = 3,
3611   .dest_bit_sizes = 0x0,
3612   .bit_size_src = -1,
3613   .num_indices = 1,
3614   .indices = {
3615      NIR_INTRINSIC_INTERP_MODE,
3616   },
3617   .index_map = {
3618      [NIR_INTRINSIC_INTERP_MODE] = 1,
3619    },
3620   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3621},
3622{
3623   .name = "load_barycentric_pixel",
3624   .num_srcs = 0,
3625   .has_dest = true,
3626   .dest_components = 2,
3627   .dest_bit_sizes = 0x0,
3628   .bit_size_src = -1,
3629   .num_indices = 1,
3630   .indices = {
3631      NIR_INTRINSIC_INTERP_MODE,
3632   },
3633   .index_map = {
3634      [NIR_INTRINSIC_INTERP_MODE] = 1,
3635    },
3636   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3637},
3638{
3639   .name = "load_barycentric_sample",
3640   .num_srcs = 0,
3641   .has_dest = true,
3642   .dest_components = 2,
3643   .dest_bit_sizes = 0x0,
3644   .bit_size_src = -1,
3645   .num_indices = 1,
3646   .indices = {
3647      NIR_INTRINSIC_INTERP_MODE,
3648   },
3649   .index_map = {
3650      [NIR_INTRINSIC_INTERP_MODE] = 1,
3651    },
3652   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3653},
3654{
3655   .name = "load_base_global_invocation_id",
3656   .num_srcs = 0,
3657   .has_dest = true,
3658   .dest_components = 3,
3659   .dest_bit_sizes = 0x60,
3660   .bit_size_src = -1,
3661   .num_indices = 0,
3662   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3663},
3664{
3665   .name = "load_base_instance",
3666   .num_srcs = 0,
3667   .has_dest = true,
3668   .dest_components = 1,
3669   .dest_bit_sizes = 0x20,
3670   .bit_size_src = -1,
3671   .num_indices = 0,
3672   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3673},
3674{
3675   .name = "load_base_vertex",
3676   .num_srcs = 0,
3677   .has_dest = true,
3678   .dest_components = 1,
3679   .dest_bit_sizes = 0x20,
3680   .bit_size_src = -1,
3681   .num_indices = 0,
3682   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3683},
3684{
3685   .name = "load_base_workgroup_id",
3686   .num_srcs = 0,
3687   .has_dest = true,
3688   .dest_components = 3,
3689   .dest_bit_sizes = 0x60,
3690   .bit_size_src = -1,
3691   .num_indices = 0,
3692   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3693},
3694{
3695   .name = "load_blend_const_color_a_float",
3696   .num_srcs = 0,
3697   .has_dest = true,
3698   .dest_components = 1,
3699   .dest_bit_sizes = 0x20,
3700   .bit_size_src = -1,
3701   .num_indices = 0,
3702   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3703},
3704{
3705   .name = "load_blend_const_color_aaaa8888_unorm",
3706   .num_srcs = 0,
3707   .has_dest = true,
3708   .dest_components = 1,
3709   .dest_bit_sizes = 0x20,
3710   .bit_size_src = -1,
3711   .num_indices = 0,
3712   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3713},
3714{
3715   .name = "load_blend_const_color_b_float",
3716   .num_srcs = 0,
3717   .has_dest = true,
3718   .dest_components = 1,
3719   .dest_bit_sizes = 0x20,
3720   .bit_size_src = -1,
3721   .num_indices = 0,
3722   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3723},
3724{
3725   .name = "load_blend_const_color_g_float",
3726   .num_srcs = 0,
3727   .has_dest = true,
3728   .dest_components = 1,
3729   .dest_bit_sizes = 0x20,
3730   .bit_size_src = -1,
3731   .num_indices = 0,
3732   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3733},
3734{
3735   .name = "load_blend_const_color_r_float",
3736   .num_srcs = 0,
3737   .has_dest = true,
3738   .dest_components = 1,
3739   .dest_bit_sizes = 0x20,
3740   .bit_size_src = -1,
3741   .num_indices = 0,
3742   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3743},
3744{
3745   .name = "load_blend_const_color_rgba",
3746   .num_srcs = 0,
3747   .has_dest = true,
3748   .dest_components = 4,
3749   .dest_bit_sizes = 0x20,
3750   .bit_size_src = -1,
3751   .num_indices = 0,
3752   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3753},
3754{
3755   .name = "load_blend_const_color_rgba8888_unorm",
3756   .num_srcs = 0,
3757   .has_dest = true,
3758   .dest_components = 1,
3759   .dest_bit_sizes = 0x20,
3760   .bit_size_src = -1,
3761   .num_indices = 0,
3762   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3763},
3764{
3765   .name = "load_btd_dss_id_intel",
3766   .num_srcs = 0,
3767   .has_dest = true,
3768   .dest_components = 1,
3769   .dest_bit_sizes = 0x20,
3770   .bit_size_src = -1,
3771   .num_indices = 0,
3772   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3773},
3774{
3775   .name = "load_btd_global_arg_addr_intel",
3776   .num_srcs = 0,
3777   .has_dest = true,
3778   .dest_components = 1,
3779   .dest_bit_sizes = 0x40,
3780   .bit_size_src = -1,
3781   .num_indices = 0,
3782   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3783},
3784{
3785   .name = "load_btd_local_arg_addr_intel",
3786   .num_srcs = 0,
3787   .has_dest = true,
3788   .dest_components = 1,
3789   .dest_bit_sizes = 0x40,
3790   .bit_size_src = -1,
3791   .num_indices = 0,
3792   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3793},
3794{
3795   .name = "load_btd_resume_sbt_addr_intel",
3796   .num_srcs = 0,
3797   .has_dest = true,
3798   .dest_components = 1,
3799   .dest_bit_sizes = 0x40,
3800   .bit_size_src = -1,
3801   .num_indices = 0,
3802   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3803},
3804{
3805   .name = "load_btd_stack_id_intel",
3806   .num_srcs = 0,
3807   .has_dest = true,
3808   .dest_components = 1,
3809   .dest_bit_sizes = 0x20,
3810   .bit_size_src = -1,
3811   .num_indices = 0,
3812   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3813},
3814{
3815   .name = "load_buffer_amd",
3816   .num_srcs = 3,
3817   .src_components = {
3818      4, 1, 1
3819   },
3820   .has_dest = true,
3821   .dest_components = 0,
3822   .dest_bit_sizes = 0x0,
3823   .bit_size_src = -1,
3824   .num_indices = 4,
3825   .indices = {
3826      NIR_INTRINSIC_BASE,
3827      NIR_INTRINSIC_IS_SWIZZLED,
3828      NIR_INTRINSIC_SLC_AMD,
3829      NIR_INTRINSIC_MEMORY_MODES,
3830   },
3831   .index_map = {
3832      [NIR_INTRINSIC_BASE] = 1,
3833      [NIR_INTRINSIC_IS_SWIZZLED] = 2,
3834      [NIR_INTRINSIC_SLC_AMD] = 3,
3835      [NIR_INTRINSIC_MEMORY_MODES] = 4,
3836    },
3837   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
3838},
3839{
3840   .name = "load_callable_sbt_addr_intel",
3841   .num_srcs = 0,
3842   .has_dest = true,
3843   .dest_components = 1,
3844   .dest_bit_sizes = 0x40,
3845   .bit_size_src = -1,
3846   .num_indices = 0,
3847   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3848},
3849{
3850   .name = "load_callable_sbt_stride_intel",
3851   .num_srcs = 0,
3852   .has_dest = true,
3853   .dest_components = 1,
3854   .dest_bit_sizes = 0x10,
3855   .bit_size_src = -1,
3856   .num_indices = 0,
3857   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3858},
3859{
3860   .name = "load_color0",
3861   .num_srcs = 0,
3862   .has_dest = true,
3863   .dest_components = 4,
3864   .dest_bit_sizes = 0x20,
3865   .bit_size_src = -1,
3866   .num_indices = 0,
3867   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3868},
3869{
3870   .name = "load_color1",
3871   .num_srcs = 0,
3872   .has_dest = true,
3873   .dest_components = 4,
3874   .dest_bit_sizes = 0x20,
3875   .bit_size_src = -1,
3876   .num_indices = 0,
3877   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3878},
3879{
3880   .name = "load_constant",
3881   .num_srcs = 1,
3882   .src_components = {
3883      1
3884   },
3885   .has_dest = true,
3886   .dest_components = 0,
3887   .dest_bit_sizes = 0x0,
3888   .bit_size_src = -1,
3889   .num_indices = 4,
3890   .indices = {
3891      NIR_INTRINSIC_BASE,
3892      NIR_INTRINSIC_RANGE,
3893      NIR_INTRINSIC_ALIGN_MUL,
3894      NIR_INTRINSIC_ALIGN_OFFSET,
3895   },
3896   .index_map = {
3897      [NIR_INTRINSIC_BASE] = 1,
3898      [NIR_INTRINSIC_RANGE] = 2,
3899      [NIR_INTRINSIC_ALIGN_MUL] = 3,
3900      [NIR_INTRINSIC_ALIGN_OFFSET] = 4,
3901    },
3902   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3903},
3904{
3905   .name = "load_constant_base_ptr",
3906   .num_srcs = 0,
3907   .has_dest = true,
3908   .dest_components = 0,
3909   .dest_bit_sizes = 0x60,
3910   .bit_size_src = -1,
3911   .num_indices = 0,
3912   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3913},
3914{
3915   .name = "load_cull_any_enabled_amd",
3916   .num_srcs = 0,
3917   .has_dest = true,
3918   .dest_components = 1,
3919   .dest_bit_sizes = 0x1,
3920   .bit_size_src = -1,
3921   .num_indices = 0,
3922   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
3923},
3924{
3925   .name = "load_cull_back_face_enabled_amd",
3926   .num_srcs = 0,
3927   .has_dest = true,
3928   .dest_components = 1,
3929   .dest_bit_sizes = 0x1,
3930   .bit_size_src = -1,
3931   .num_indices = 0,
3932   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
3933},
3934{
3935   .name = "load_cull_ccw_amd",
3936   .num_srcs = 0,
3937   .has_dest = true,
3938   .dest_components = 1,
3939   .dest_bit_sizes = 0x1,
3940   .bit_size_src = -1,
3941   .num_indices = 0,
3942   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
3943},
3944{
3945   .name = "load_cull_front_face_enabled_amd",
3946   .num_srcs = 0,
3947   .has_dest = true,
3948   .dest_components = 1,
3949   .dest_bit_sizes = 0x1,
3950   .bit_size_src = -1,
3951   .num_indices = 0,
3952   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
3953},
3954{
3955   .name = "load_cull_small_prim_precision_amd",
3956   .num_srcs = 0,
3957   .has_dest = true,
3958   .dest_components = 1,
3959   .dest_bit_sizes = 0x20,
3960   .bit_size_src = -1,
3961   .num_indices = 0,
3962   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
3963},
3964{
3965   .name = "load_cull_small_primitives_enabled_amd",
3966   .num_srcs = 0,
3967   .has_dest = true,
3968   .dest_components = 1,
3969   .dest_bit_sizes = 0x1,
3970   .bit_size_src = -1,
3971   .num_indices = 0,
3972   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
3973},
3974{
3975   .name = "load_deref",
3976   .num_srcs = 1,
3977   .src_components = {
3978      -1
3979   },
3980   .has_dest = true,
3981   .dest_components = 0,
3982   .dest_bit_sizes = 0x0,
3983   .bit_size_src = -1,
3984   .num_indices = 1,
3985   .indices = {
3986      NIR_INTRINSIC_ACCESS,
3987   },
3988   .index_map = {
3989      [NIR_INTRINSIC_ACCESS] = 1,
3990    },
3991   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
3992},
3993{
3994   .name = "load_deref_block_intel",
3995   .num_srcs = 1,
3996   .src_components = {
3997      -1
3998   },
3999   .has_dest = true,
4000   .dest_components = 0,
4001   .dest_bit_sizes = 0x0,
4002   .bit_size_src = -1,
4003   .num_indices = 1,
4004   .indices = {
4005      NIR_INTRINSIC_ACCESS,
4006   },
4007   .index_map = {
4008      [NIR_INTRINSIC_ACCESS] = 1,
4009    },
4010   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
4011},
4012{
4013   .name = "load_desc_set_address_intel",
4014   .num_srcs = 1,
4015   .src_components = {
4016      1
4017   },
4018   .has_dest = true,
4019   .dest_components = 1,
4020   .dest_bit_sizes = 0x40,
4021   .bit_size_src = -1,
4022   .num_indices = 0,
4023   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4024},
4025{
4026   .name = "load_draw_id",
4027   .num_srcs = 0,
4028   .has_dest = true,
4029   .dest_components = 1,
4030   .dest_bit_sizes = 0x20,
4031   .bit_size_src = -1,
4032   .num_indices = 0,
4033   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4034},
4035{
4036   .name = "load_fb_layers_v3d",
4037   .num_srcs = 0,
4038   .has_dest = true,
4039   .dest_components = 1,
4040   .dest_bit_sizes = 0x0,
4041   .bit_size_src = -1,
4042   .num_indices = 0,
4043   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4044},
4045{
4046   .name = "load_first_vertex",
4047   .num_srcs = 0,
4048   .has_dest = true,
4049   .dest_components = 1,
4050   .dest_bit_sizes = 0x20,
4051   .bit_size_src = -1,
4052   .num_indices = 0,
4053   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4054},
4055{
4056   .name = "load_frag_coord",
4057   .num_srcs = 0,
4058   .has_dest = true,
4059   .dest_components = 4,
4060   .dest_bit_sizes = 0x20,
4061   .bit_size_src = -1,
4062   .num_indices = 0,
4063   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4064},
4065{
4066   .name = "load_frag_shading_rate",
4067   .num_srcs = 0,
4068   .has_dest = true,
4069   .dest_components = 1,
4070   .dest_bit_sizes = 0x20,
4071   .bit_size_src = -1,
4072   .num_indices = 0,
4073   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4074},
4075{
4076   .name = "load_front_face",
4077   .num_srcs = 0,
4078   .has_dest = true,
4079   .dest_components = 1,
4080   .dest_bit_sizes = 0x21,
4081   .bit_size_src = -1,
4082   .num_indices = 0,
4083   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4084},
4085{
4086   .name = "load_fs_input_interp_deltas",
4087   .num_srcs = 1,
4088   .src_components = {
4089      1
4090   },
4091   .has_dest = true,
4092   .dest_components = 3,
4093   .dest_bit_sizes = 0x0,
4094   .bit_size_src = -1,
4095   .num_indices = 3,
4096   .indices = {
4097      NIR_INTRINSIC_BASE,
4098      NIR_INTRINSIC_COMPONENT,
4099      NIR_INTRINSIC_IO_SEMANTICS,
4100   },
4101   .index_map = {
4102      [NIR_INTRINSIC_BASE] = 1,
4103      [NIR_INTRINSIC_COMPONENT] = 2,
4104      [NIR_INTRINSIC_IO_SEMANTICS] = 3,
4105    },
4106   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4107},
4108{
4109   .name = "load_global",
4110   .num_srcs = 1,
4111   .src_components = {
4112      1
4113   },
4114   .has_dest = true,
4115   .dest_components = 0,
4116   .dest_bit_sizes = 0x0,
4117   .bit_size_src = -1,
4118   .num_indices = 3,
4119   .indices = {
4120      NIR_INTRINSIC_ACCESS,
4121      NIR_INTRINSIC_ALIGN_MUL,
4122      NIR_INTRINSIC_ALIGN_OFFSET,
4123   },
4124   .index_map = {
4125      [NIR_INTRINSIC_ACCESS] = 1,
4126      [NIR_INTRINSIC_ALIGN_MUL] = 2,
4127      [NIR_INTRINSIC_ALIGN_OFFSET] = 3,
4128    },
4129   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
4130},
4131{
4132   .name = "load_global_block_intel",
4133   .num_srcs = 1,
4134   .src_components = {
4135      1
4136   },
4137   .has_dest = true,
4138   .dest_components = 0,
4139   .dest_bit_sizes = 0x0,
4140   .bit_size_src = -1,
4141   .num_indices = 3,
4142   .indices = {
4143      NIR_INTRINSIC_ACCESS,
4144      NIR_INTRINSIC_ALIGN_MUL,
4145      NIR_INTRINSIC_ALIGN_OFFSET,
4146   },
4147   .index_map = {
4148      [NIR_INTRINSIC_ACCESS] = 1,
4149      [NIR_INTRINSIC_ALIGN_MUL] = 2,
4150      [NIR_INTRINSIC_ALIGN_OFFSET] = 3,
4151    },
4152   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
4153},
4154{
4155   .name = "load_global_const_block_intel",
4156   .num_srcs = 2,
4157   .src_components = {
4158      1, 1
4159   },
4160   .has_dest = true,
4161   .dest_components = 0,
4162   .dest_bit_sizes = 0x20,
4163   .bit_size_src = -1,
4164   .num_indices = 1,
4165   .indices = {
4166      NIR_INTRINSIC_BASE,
4167   },
4168   .index_map = {
4169      [NIR_INTRINSIC_BASE] = 1,
4170    },
4171   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4172},
4173{
4174   .name = "load_global_constant",
4175   .num_srcs = 1,
4176   .src_components = {
4177      1
4178   },
4179   .has_dest = true,
4180   .dest_components = 0,
4181   .dest_bit_sizes = 0x0,
4182   .bit_size_src = -1,
4183   .num_indices = 3,
4184   .indices = {
4185      NIR_INTRINSIC_ACCESS,
4186      NIR_INTRINSIC_ALIGN_MUL,
4187      NIR_INTRINSIC_ALIGN_OFFSET,
4188   },
4189   .index_map = {
4190      [NIR_INTRINSIC_ACCESS] = 1,
4191      [NIR_INTRINSIC_ALIGN_MUL] = 2,
4192      [NIR_INTRINSIC_ALIGN_OFFSET] = 3,
4193    },
4194   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4195},
4196{
4197   .name = "load_global_constant_bounded",
4198   .num_srcs = 3,
4199   .src_components = {
4200      1, 1, 1
4201   },
4202   .has_dest = true,
4203   .dest_components = 0,
4204   .dest_bit_sizes = 0x0,
4205   .bit_size_src = -1,
4206   .num_indices = 3,
4207   .indices = {
4208      NIR_INTRINSIC_ACCESS,
4209      NIR_INTRINSIC_ALIGN_MUL,
4210      NIR_INTRINSIC_ALIGN_OFFSET,
4211   },
4212   .index_map = {
4213      [NIR_INTRINSIC_ACCESS] = 1,
4214      [NIR_INTRINSIC_ALIGN_MUL] = 2,
4215      [NIR_INTRINSIC_ALIGN_OFFSET] = 3,
4216    },
4217   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4218},
4219{
4220   .name = "load_global_constant_offset",
4221   .num_srcs = 2,
4222   .src_components = {
4223      1, 1
4224   },
4225   .has_dest = true,
4226   .dest_components = 0,
4227   .dest_bit_sizes = 0x0,
4228   .bit_size_src = -1,
4229   .num_indices = 3,
4230   .indices = {
4231      NIR_INTRINSIC_ACCESS,
4232      NIR_INTRINSIC_ALIGN_MUL,
4233      NIR_INTRINSIC_ALIGN_OFFSET,
4234   },
4235   .index_map = {
4236      [NIR_INTRINSIC_ACCESS] = 1,
4237      [NIR_INTRINSIC_ALIGN_MUL] = 2,
4238      [NIR_INTRINSIC_ALIGN_OFFSET] = 3,
4239    },
4240   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4241},
4242{
4243   .name = "load_global_invocation_id",
4244   .num_srcs = 0,
4245   .has_dest = true,
4246   .dest_components = 3,
4247   .dest_bit_sizes = 0x60,
4248   .bit_size_src = -1,
4249   .num_indices = 0,
4250   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4251},
4252{
4253   .name = "load_global_invocation_id_zero_base",
4254   .num_srcs = 0,
4255   .has_dest = true,
4256   .dest_components = 3,
4257   .dest_bit_sizes = 0x60,
4258   .bit_size_src = -1,
4259   .num_indices = 0,
4260   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4261},
4262{
4263   .name = "load_global_invocation_index",
4264   .num_srcs = 0,
4265   .has_dest = true,
4266   .dest_components = 1,
4267   .dest_bit_sizes = 0x60,
4268   .bit_size_src = -1,
4269   .num_indices = 0,
4270   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4271},
4272{
4273   .name = "load_global_ir3",
4274   .num_srcs = 2,
4275   .src_components = {
4276      2, 1
4277   },
4278   .has_dest = true,
4279   .dest_components = 0,
4280   .dest_bit_sizes = 0x0,
4281   .bit_size_src = -1,
4282   .num_indices = 3,
4283   .indices = {
4284      NIR_INTRINSIC_ACCESS,
4285      NIR_INTRINSIC_ALIGN_MUL,
4286      NIR_INTRINSIC_ALIGN_OFFSET,
4287   },
4288   .index_map = {
4289      [NIR_INTRINSIC_ACCESS] = 1,
4290      [NIR_INTRINSIC_ALIGN_MUL] = 2,
4291      [NIR_INTRINSIC_ALIGN_OFFSET] = 3,
4292    },
4293   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
4294},
4295{
4296   .name = "load_gs_header_ir3",
4297   .num_srcs = 0,
4298   .has_dest = true,
4299   .dest_components = 1,
4300   .dest_bit_sizes = 0x20,
4301   .bit_size_src = -1,
4302   .num_indices = 0,
4303   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4304},
4305{
4306   .name = "load_gs_vertex_offset_amd",
4307   .num_srcs = 0,
4308   .has_dest = true,
4309   .dest_components = 1,
4310   .dest_bit_sizes = 0x20,
4311   .bit_size_src = -1,
4312   .num_indices = 1,
4313   .indices = {
4314      NIR_INTRINSIC_BASE,
4315   },
4316   .index_map = {
4317      [NIR_INTRINSIC_BASE] = 1,
4318    },
4319   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4320},
4321{
4322   .name = "load_helper_invocation",
4323   .num_srcs = 0,
4324   .has_dest = true,
4325   .dest_components = 1,
4326   .dest_bit_sizes = 0x21,
4327   .bit_size_src = -1,
4328   .num_indices = 0,
4329   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4330},
4331{
4332   .name = "load_hs_patch_stride_ir3",
4333   .num_srcs = 0,
4334   .has_dest = true,
4335   .dest_components = 1,
4336   .dest_bit_sizes = 0x20,
4337   .bit_size_src = -1,
4338   .num_indices = 0,
4339   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4340},
4341{
4342   .name = "load_initial_edgeflags_amd",
4343   .num_srcs = 0,
4344   .has_dest = true,
4345   .dest_components = 1,
4346   .dest_bit_sizes = 0x20,
4347   .bit_size_src = -1,
4348   .num_indices = 0,
4349   .flags = 0,
4350},
4351{
4352   .name = "load_input",
4353   .num_srcs = 1,
4354   .src_components = {
4355      1
4356   },
4357   .has_dest = true,
4358   .dest_components = 0,
4359   .dest_bit_sizes = 0x0,
4360   .bit_size_src = -1,
4361   .num_indices = 4,
4362   .indices = {
4363      NIR_INTRINSIC_BASE,
4364      NIR_INTRINSIC_COMPONENT,
4365      NIR_INTRINSIC_DEST_TYPE,
4366      NIR_INTRINSIC_IO_SEMANTICS,
4367   },
4368   .index_map = {
4369      [NIR_INTRINSIC_BASE] = 1,
4370      [NIR_INTRINSIC_COMPONENT] = 2,
4371      [NIR_INTRINSIC_DEST_TYPE] = 3,
4372      [NIR_INTRINSIC_IO_SEMANTICS] = 4,
4373    },
4374   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4375},
4376{
4377   .name = "load_input_vertex",
4378   .num_srcs = 2,
4379   .src_components = {
4380      1, 1
4381   },
4382   .has_dest = true,
4383   .dest_components = 0,
4384   .dest_bit_sizes = 0x0,
4385   .bit_size_src = -1,
4386   .num_indices = 4,
4387   .indices = {
4388      NIR_INTRINSIC_BASE,
4389      NIR_INTRINSIC_COMPONENT,
4390      NIR_INTRINSIC_DEST_TYPE,
4391      NIR_INTRINSIC_IO_SEMANTICS,
4392   },
4393   .index_map = {
4394      [NIR_INTRINSIC_BASE] = 1,
4395      [NIR_INTRINSIC_COMPONENT] = 2,
4396      [NIR_INTRINSIC_DEST_TYPE] = 3,
4397      [NIR_INTRINSIC_IO_SEMANTICS] = 4,
4398    },
4399   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4400},
4401{
4402   .name = "load_instance_id",
4403   .num_srcs = 0,
4404   .has_dest = true,
4405   .dest_components = 1,
4406   .dest_bit_sizes = 0x20,
4407   .bit_size_src = -1,
4408   .num_indices = 0,
4409   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4410},
4411{
4412   .name = "load_interpolated_input",
4413   .num_srcs = 2,
4414   .src_components = {
4415      2, 1
4416   },
4417   .has_dest = true,
4418   .dest_components = 0,
4419   .dest_bit_sizes = 0x0,
4420   .bit_size_src = -1,
4421   .num_indices = 4,
4422   .indices = {
4423      NIR_INTRINSIC_BASE,
4424      NIR_INTRINSIC_COMPONENT,
4425      NIR_INTRINSIC_DEST_TYPE,
4426      NIR_INTRINSIC_IO_SEMANTICS,
4427   },
4428   .index_map = {
4429      [NIR_INTRINSIC_BASE] = 1,
4430      [NIR_INTRINSIC_COMPONENT] = 2,
4431      [NIR_INTRINSIC_DEST_TYPE] = 3,
4432      [NIR_INTRINSIC_IO_SEMANTICS] = 4,
4433    },
4434   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4435},
4436{
4437   .name = "load_intersection_opaque_amd",
4438   .num_srcs = 0,
4439   .has_dest = true,
4440   .dest_components = 1,
4441   .dest_bit_sizes = 0x1,
4442   .bit_size_src = -1,
4443   .num_indices = 0,
4444   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4445},
4446{
4447   .name = "load_invocation_id",
4448   .num_srcs = 0,
4449   .has_dest = true,
4450   .dest_components = 1,
4451   .dest_bit_sizes = 0x20,
4452   .bit_size_src = -1,
4453   .num_indices = 0,
4454   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4455},
4456{
4457   .name = "load_is_indexed_draw",
4458   .num_srcs = 0,
4459   .has_dest = true,
4460   .dest_components = 1,
4461   .dest_bit_sizes = 0x20,
4462   .bit_size_src = -1,
4463   .num_indices = 0,
4464   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4465},
4466{
4467   .name = "load_kernel_input",
4468   .num_srcs = 1,
4469   .src_components = {
4470      1
4471   },
4472   .has_dest = true,
4473   .dest_components = 0,
4474   .dest_bit_sizes = 0x0,
4475   .bit_size_src = -1,
4476   .num_indices = 4,
4477   .indices = {
4478      NIR_INTRINSIC_BASE,
4479      NIR_INTRINSIC_RANGE,
4480      NIR_INTRINSIC_ALIGN_MUL,
4481      NIR_INTRINSIC_ALIGN_OFFSET,
4482   },
4483   .index_map = {
4484      [NIR_INTRINSIC_BASE] = 1,
4485      [NIR_INTRINSIC_RANGE] = 2,
4486      [NIR_INTRINSIC_ALIGN_MUL] = 3,
4487      [NIR_INTRINSIC_ALIGN_OFFSET] = 4,
4488    },
4489   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4490},
4491{
4492   .name = "load_layer_id",
4493   .num_srcs = 0,
4494   .has_dest = true,
4495   .dest_components = 1,
4496   .dest_bit_sizes = 0x20,
4497   .bit_size_src = -1,
4498   .num_indices = 0,
4499   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4500},
4501{
4502   .name = "load_leaf_opaque_intel",
4503   .num_srcs = 0,
4504   .has_dest = true,
4505   .dest_components = 1,
4506   .dest_bit_sizes = 0x1,
4507   .bit_size_src = -1,
4508   .num_indices = 0,
4509   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4510},
4511{
4512   .name = "load_leaf_procedural_intel",
4513   .num_srcs = 0,
4514   .has_dest = true,
4515   .dest_components = 1,
4516   .dest_bit_sizes = 0x1,
4517   .bit_size_src = -1,
4518   .num_indices = 0,
4519   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4520},
4521{
4522   .name = "load_line_coord",
4523   .num_srcs = 0,
4524   .has_dest = true,
4525   .dest_components = 1,
4526   .dest_bit_sizes = 0x20,
4527   .bit_size_src = -1,
4528   .num_indices = 0,
4529   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4530},
4531{
4532   .name = "load_line_width",
4533   .num_srcs = 0,
4534   .has_dest = true,
4535   .dest_components = 1,
4536   .dest_bit_sizes = 0x20,
4537   .bit_size_src = -1,
4538   .num_indices = 0,
4539   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4540},
4541{
4542   .name = "load_local_invocation_id",
4543   .num_srcs = 0,
4544   .has_dest = true,
4545   .dest_components = 3,
4546   .dest_bit_sizes = 0x20,
4547   .bit_size_src = -1,
4548   .num_indices = 0,
4549   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4550},
4551{
4552   .name = "load_local_invocation_index",
4553   .num_srcs = 0,
4554   .has_dest = true,
4555   .dest_components = 1,
4556   .dest_bit_sizes = 0x20,
4557   .bit_size_src = -1,
4558   .num_indices = 0,
4559   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4560},
4561{
4562   .name = "load_local_shared_r600",
4563   .num_srcs = 1,
4564   .src_components = {
4565      0
4566   },
4567   .has_dest = true,
4568   .dest_components = 0,
4569   .dest_bit_sizes = 0x0,
4570   .bit_size_src = -1,
4571   .num_indices = 0,
4572   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
4573},
4574{
4575   .name = "load_num_subgroups",
4576   .num_srcs = 0,
4577   .has_dest = true,
4578   .dest_components = 1,
4579   .dest_bit_sizes = 0x20,
4580   .bit_size_src = -1,
4581   .num_indices = 0,
4582   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4583},
4584{
4585   .name = "load_num_workgroups",
4586   .num_srcs = 0,
4587   .has_dest = true,
4588   .dest_components = 3,
4589   .dest_bit_sizes = 0x60,
4590   .bit_size_src = -1,
4591   .num_indices = 0,
4592   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4593},
4594{
4595   .name = "load_output",
4596   .num_srcs = 1,
4597   .src_components = {
4598      1
4599   },
4600   .has_dest = true,
4601   .dest_components = 0,
4602   .dest_bit_sizes = 0x0,
4603   .bit_size_src = -1,
4604   .num_indices = 4,
4605   .indices = {
4606      NIR_INTRINSIC_BASE,
4607      NIR_INTRINSIC_COMPONENT,
4608      NIR_INTRINSIC_DEST_TYPE,
4609      NIR_INTRINSIC_IO_SEMANTICS,
4610   },
4611   .index_map = {
4612      [NIR_INTRINSIC_BASE] = 1,
4613      [NIR_INTRINSIC_COMPONENT] = 2,
4614      [NIR_INTRINSIC_DEST_TYPE] = 3,
4615      [NIR_INTRINSIC_IO_SEMANTICS] = 4,
4616    },
4617   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
4618},
4619{
4620   .name = "load_packed_passthrough_primitive_amd",
4621   .num_srcs = 0,
4622   .has_dest = true,
4623   .dest_components = 1,
4624   .dest_bit_sizes = 0x20,
4625   .bit_size_src = -1,
4626   .num_indices = 0,
4627   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4628},
4629{
4630   .name = "load_param",
4631   .num_srcs = 0,
4632   .has_dest = true,
4633   .dest_components = 0,
4634   .dest_bit_sizes = 0x0,
4635   .bit_size_src = -1,
4636   .num_indices = 1,
4637   .indices = {
4638      NIR_INTRINSIC_PARAM_IDX,
4639   },
4640   .index_map = {
4641      [NIR_INTRINSIC_PARAM_IDX] = 1,
4642    },
4643   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
4644},
4645{
4646   .name = "load_patch_vertices_in",
4647   .num_srcs = 0,
4648   .has_dest = true,
4649   .dest_components = 1,
4650   .dest_bit_sizes = 0x20,
4651   .bit_size_src = -1,
4652   .num_indices = 0,
4653   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4654},
4655{
4656   .name = "load_per_primitive_output",
4657   .num_srcs = 2,
4658   .src_components = {
4659      1, 1
4660   },
4661   .has_dest = true,
4662   .dest_components = 0,
4663   .dest_bit_sizes = 0x0,
4664   .bit_size_src = -1,
4665   .num_indices = 4,
4666   .indices = {
4667      NIR_INTRINSIC_BASE,
4668      NIR_INTRINSIC_COMPONENT,
4669      NIR_INTRINSIC_DEST_TYPE,
4670      NIR_INTRINSIC_IO_SEMANTICS,
4671   },
4672   .index_map = {
4673      [NIR_INTRINSIC_BASE] = 1,
4674      [NIR_INTRINSIC_COMPONENT] = 2,
4675      [NIR_INTRINSIC_DEST_TYPE] = 3,
4676      [NIR_INTRINSIC_IO_SEMANTICS] = 4,
4677    },
4678   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
4679},
4680{
4681   .name = "load_per_vertex_input",
4682   .num_srcs = 2,
4683   .src_components = {
4684      1, 1
4685   },
4686   .has_dest = true,
4687   .dest_components = 0,
4688   .dest_bit_sizes = 0x0,
4689   .bit_size_src = -1,
4690   .num_indices = 4,
4691   .indices = {
4692      NIR_INTRINSIC_BASE,
4693      NIR_INTRINSIC_COMPONENT,
4694      NIR_INTRINSIC_DEST_TYPE,
4695      NIR_INTRINSIC_IO_SEMANTICS,
4696   },
4697   .index_map = {
4698      [NIR_INTRINSIC_BASE] = 1,
4699      [NIR_INTRINSIC_COMPONENT] = 2,
4700      [NIR_INTRINSIC_DEST_TYPE] = 3,
4701      [NIR_INTRINSIC_IO_SEMANTICS] = 4,
4702    },
4703   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4704},
4705{
4706   .name = "load_per_vertex_output",
4707   .num_srcs = 2,
4708   .src_components = {
4709      1, 1
4710   },
4711   .has_dest = true,
4712   .dest_components = 0,
4713   .dest_bit_sizes = 0x0,
4714   .bit_size_src = -1,
4715   .num_indices = 4,
4716   .indices = {
4717      NIR_INTRINSIC_BASE,
4718      NIR_INTRINSIC_COMPONENT,
4719      NIR_INTRINSIC_DEST_TYPE,
4720      NIR_INTRINSIC_IO_SEMANTICS,
4721   },
4722   .index_map = {
4723      [NIR_INTRINSIC_BASE] = 1,
4724      [NIR_INTRINSIC_COMPONENT] = 2,
4725      [NIR_INTRINSIC_DEST_TYPE] = 3,
4726      [NIR_INTRINSIC_IO_SEMANTICS] = 4,
4727    },
4728   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
4729},
4730{
4731   .name = "load_point_coord",
4732   .num_srcs = 0,
4733   .has_dest = true,
4734   .dest_components = 2,
4735   .dest_bit_sizes = 0x20,
4736   .bit_size_src = -1,
4737   .num_indices = 0,
4738   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4739},
4740{
4741   .name = "load_primitive_id",
4742   .num_srcs = 0,
4743   .has_dest = true,
4744   .dest_components = 1,
4745   .dest_bit_sizes = 0x20,
4746   .bit_size_src = -1,
4747   .num_indices = 0,
4748   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4749},
4750{
4751   .name = "load_primitive_location_ir3",
4752   .num_srcs = 0,
4753   .has_dest = true,
4754   .dest_components = 1,
4755   .dest_bit_sizes = 0x20,
4756   .bit_size_src = -1,
4757   .num_indices = 1,
4758   .indices = {
4759      NIR_INTRINSIC_DRIVER_LOCATION,
4760   },
4761   .index_map = {
4762      [NIR_INTRINSIC_DRIVER_LOCATION] = 1,
4763    },
4764   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4765},
4766{
4767   .name = "load_printf_buffer_address",
4768   .num_srcs = 0,
4769   .has_dest = true,
4770   .dest_components = 1,
4771   .dest_bit_sizes = 0x60,
4772   .bit_size_src = -1,
4773   .num_indices = 0,
4774   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4775},
4776{
4777   .name = "load_ptr_dxil",
4778   .num_srcs = 2,
4779   .src_components = {
4780      1, 1
4781   },
4782   .has_dest = true,
4783   .dest_components = 0,
4784   .dest_bit_sizes = 0x0,
4785   .bit_size_src = -1,
4786   .num_indices = 0,
4787   .flags = 0,
4788},
4789{
4790   .name = "load_push_constant",
4791   .num_srcs = 1,
4792   .src_components = {
4793      1
4794   },
4795   .has_dest = true,
4796   .dest_components = 0,
4797   .dest_bit_sizes = 0x0,
4798   .bit_size_src = -1,
4799   .num_indices = 2,
4800   .indices = {
4801      NIR_INTRINSIC_BASE,
4802      NIR_INTRINSIC_RANGE,
4803   },
4804   .index_map = {
4805      [NIR_INTRINSIC_BASE] = 1,
4806      [NIR_INTRINSIC_RANGE] = 2,
4807    },
4808   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4809},
4810{
4811   .name = "load_raw_output_pan",
4812   .num_srcs = 1,
4813   .src_components = {
4814      1
4815   },
4816   .has_dest = true,
4817   .dest_components = 0,
4818   .dest_bit_sizes = 0x0,
4819   .bit_size_src = -1,
4820   .num_indices = 1,
4821   .indices = {
4822      NIR_INTRINSIC_BASE,
4823   },
4824   .index_map = {
4825      [NIR_INTRINSIC_BASE] = 1,
4826    },
4827   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4828},
4829{
4830   .name = "load_ray_base_mem_addr_intel",
4831   .num_srcs = 0,
4832   .has_dest = true,
4833   .dest_components = 1,
4834   .dest_bit_sizes = 0x40,
4835   .bit_size_src = -1,
4836   .num_indices = 0,
4837   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4838},
4839{
4840   .name = "load_ray_flags",
4841   .num_srcs = 0,
4842   .has_dest = true,
4843   .dest_components = 1,
4844   .dest_bit_sizes = 0x20,
4845   .bit_size_src = -1,
4846   .num_indices = 0,
4847   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4848},
4849{
4850   .name = "load_ray_geometry_index",
4851   .num_srcs = 0,
4852   .has_dest = true,
4853   .dest_components = 1,
4854   .dest_bit_sizes = 0x20,
4855   .bit_size_src = -1,
4856   .num_indices = 0,
4857   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4858},
4859{
4860   .name = "load_ray_hit_kind",
4861   .num_srcs = 0,
4862   .has_dest = true,
4863   .dest_components = 1,
4864   .dest_bit_sizes = 0x20,
4865   .bit_size_src = -1,
4866   .num_indices = 0,
4867   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4868},
4869{
4870   .name = "load_ray_hit_sbt_addr_intel",
4871   .num_srcs = 0,
4872   .has_dest = true,
4873   .dest_components = 1,
4874   .dest_bit_sizes = 0x40,
4875   .bit_size_src = -1,
4876   .num_indices = 0,
4877   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4878},
4879{
4880   .name = "load_ray_hit_sbt_stride_intel",
4881   .num_srcs = 0,
4882   .has_dest = true,
4883   .dest_components = 1,
4884   .dest_bit_sizes = 0x10,
4885   .bit_size_src = -1,
4886   .num_indices = 0,
4887   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4888},
4889{
4890   .name = "load_ray_hw_stack_size_intel",
4891   .num_srcs = 0,
4892   .has_dest = true,
4893   .dest_components = 1,
4894   .dest_bit_sizes = 0x20,
4895   .bit_size_src = -1,
4896   .num_indices = 0,
4897   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4898},
4899{
4900   .name = "load_ray_instance_custom_index",
4901   .num_srcs = 0,
4902   .has_dest = true,
4903   .dest_components = 1,
4904   .dest_bit_sizes = 0x20,
4905   .bit_size_src = -1,
4906   .num_indices = 0,
4907   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4908},
4909{
4910   .name = "load_ray_launch_id",
4911   .num_srcs = 0,
4912   .has_dest = true,
4913   .dest_components = 3,
4914   .dest_bit_sizes = 0x20,
4915   .bit_size_src = -1,
4916   .num_indices = 0,
4917   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4918},
4919{
4920   .name = "load_ray_launch_size",
4921   .num_srcs = 0,
4922   .has_dest = true,
4923   .dest_components = 3,
4924   .dest_bit_sizes = 0x20,
4925   .bit_size_src = -1,
4926   .num_indices = 0,
4927   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4928},
4929{
4930   .name = "load_ray_miss_sbt_addr_intel",
4931   .num_srcs = 0,
4932   .has_dest = true,
4933   .dest_components = 1,
4934   .dest_bit_sizes = 0x40,
4935   .bit_size_src = -1,
4936   .num_indices = 0,
4937   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4938},
4939{
4940   .name = "load_ray_miss_sbt_stride_intel",
4941   .num_srcs = 0,
4942   .has_dest = true,
4943   .dest_components = 1,
4944   .dest_bit_sizes = 0x10,
4945   .bit_size_src = -1,
4946   .num_indices = 0,
4947   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4948},
4949{
4950   .name = "load_ray_num_dss_rt_stacks_intel",
4951   .num_srcs = 0,
4952   .has_dest = true,
4953   .dest_components = 1,
4954   .dest_bit_sizes = 0x20,
4955   .bit_size_src = -1,
4956   .num_indices = 0,
4957   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4958},
4959{
4960   .name = "load_ray_object_direction",
4961   .num_srcs = 0,
4962   .has_dest = true,
4963   .dest_components = 3,
4964   .dest_bit_sizes = 0x20,
4965   .bit_size_src = -1,
4966   .num_indices = 0,
4967   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4968},
4969{
4970   .name = "load_ray_object_origin",
4971   .num_srcs = 0,
4972   .has_dest = true,
4973   .dest_components = 3,
4974   .dest_bit_sizes = 0x20,
4975   .bit_size_src = -1,
4976   .num_indices = 0,
4977   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4978},
4979{
4980   .name = "load_ray_object_to_world",
4981   .num_srcs = 0,
4982   .has_dest = true,
4983   .dest_components = 3,
4984   .dest_bit_sizes = 0x20,
4985   .bit_size_src = -1,
4986   .num_indices = 1,
4987   .indices = {
4988      NIR_INTRINSIC_COLUMN,
4989   },
4990   .index_map = {
4991      [NIR_INTRINSIC_COLUMN] = 1,
4992    },
4993   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
4994},
4995{
4996   .name = "load_ray_sw_stack_size_intel",
4997   .num_srcs = 0,
4998   .has_dest = true,
4999   .dest_components = 1,
5000   .dest_bit_sizes = 0x20,
5001   .bit_size_src = -1,
5002   .num_indices = 0,
5003   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5004},
5005{
5006   .name = "load_ray_t_max",
5007   .num_srcs = 0,
5008   .has_dest = true,
5009   .dest_components = 1,
5010   .dest_bit_sizes = 0x20,
5011   .bit_size_src = -1,
5012   .num_indices = 0,
5013   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5014},
5015{
5016   .name = "load_ray_t_min",
5017   .num_srcs = 0,
5018   .has_dest = true,
5019   .dest_components = 1,
5020   .dest_bit_sizes = 0x20,
5021   .bit_size_src = -1,
5022   .num_indices = 0,
5023   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5024},
5025{
5026   .name = "load_ray_world_direction",
5027   .num_srcs = 0,
5028   .has_dest = true,
5029   .dest_components = 3,
5030   .dest_bit_sizes = 0x20,
5031   .bit_size_src = -1,
5032   .num_indices = 0,
5033   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5034},
5035{
5036   .name = "load_ray_world_origin",
5037   .num_srcs = 0,
5038   .has_dest = true,
5039   .dest_components = 3,
5040   .dest_bit_sizes = 0x20,
5041   .bit_size_src = -1,
5042   .num_indices = 0,
5043   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5044},
5045{
5046   .name = "load_ray_world_to_object",
5047   .num_srcs = 0,
5048   .has_dest = true,
5049   .dest_components = 3,
5050   .dest_bit_sizes = 0x20,
5051   .bit_size_src = -1,
5052   .num_indices = 1,
5053   .indices = {
5054      NIR_INTRINSIC_COLUMN,
5055   },
5056   .index_map = {
5057      [NIR_INTRINSIC_COLUMN] = 1,
5058    },
5059   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5060},
5061{
5062   .name = "load_rel_patch_id_ir3",
5063   .num_srcs = 0,
5064   .has_dest = true,
5065   .dest_components = 1,
5066   .dest_bit_sizes = 0x20,
5067   .bit_size_src = -1,
5068   .num_indices = 0,
5069   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5070},
5071{
5072   .name = "load_reloc_const_intel",
5073   .num_srcs = 0,
5074   .has_dest = true,
5075   .dest_components = 1,
5076   .dest_bit_sizes = 0x20,
5077   .bit_size_src = -1,
5078   .num_indices = 1,
5079   .indices = {
5080      NIR_INTRINSIC_PARAM_IDX,
5081   },
5082   .index_map = {
5083      [NIR_INTRINSIC_PARAM_IDX] = 1,
5084    },
5085   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5086},
5087{
5088   .name = "load_ring_es2gs_offset_amd",
5089   .num_srcs = 0,
5090   .has_dest = true,
5091   .dest_components = 1,
5092   .dest_bit_sizes = 0x20,
5093   .bit_size_src = -1,
5094   .num_indices = 0,
5095   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5096},
5097{
5098   .name = "load_ring_esgs_amd",
5099   .num_srcs = 0,
5100   .has_dest = true,
5101   .dest_components = 4,
5102   .dest_bit_sizes = 0x20,
5103   .bit_size_src = -1,
5104   .num_indices = 0,
5105   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5106},
5107{
5108   .name = "load_ring_tess_factors_amd",
5109   .num_srcs = 0,
5110   .has_dest = true,
5111   .dest_components = 4,
5112   .dest_bit_sizes = 0x20,
5113   .bit_size_src = -1,
5114   .num_indices = 0,
5115   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5116},
5117{
5118   .name = "load_ring_tess_factors_offset_amd",
5119   .num_srcs = 0,
5120   .has_dest = true,
5121   .dest_components = 1,
5122   .dest_bit_sizes = 0x20,
5123   .bit_size_src = -1,
5124   .num_indices = 0,
5125   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5126},
5127{
5128   .name = "load_ring_tess_offchip_amd",
5129   .num_srcs = 0,
5130   .has_dest = true,
5131   .dest_components = 4,
5132   .dest_bit_sizes = 0x20,
5133   .bit_size_src = -1,
5134   .num_indices = 0,
5135   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5136},
5137{
5138   .name = "load_ring_tess_offchip_offset_amd",
5139   .num_srcs = 0,
5140   .has_dest = true,
5141   .dest_components = 1,
5142   .dest_bit_sizes = 0x20,
5143   .bit_size_src = -1,
5144   .num_indices = 0,
5145   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5146},
5147{
5148   .name = "load_rt_arg_scratch_offset_amd",
5149   .num_srcs = 0,
5150   .has_dest = true,
5151   .dest_components = 1,
5152   .dest_bit_sizes = 0x20,
5153   .bit_size_src = -1,
5154   .num_indices = 0,
5155   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5156},
5157{
5158   .name = "load_sample_id",
5159   .num_srcs = 0,
5160   .has_dest = true,
5161   .dest_components = 1,
5162   .dest_bit_sizes = 0x20,
5163   .bit_size_src = -1,
5164   .num_indices = 0,
5165   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5166},
5167{
5168   .name = "load_sample_id_no_per_sample",
5169   .num_srcs = 0,
5170   .has_dest = true,
5171   .dest_components = 1,
5172   .dest_bit_sizes = 0x20,
5173   .bit_size_src = -1,
5174   .num_indices = 0,
5175   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5176},
5177{
5178   .name = "load_sample_mask_in",
5179   .num_srcs = 0,
5180   .has_dest = true,
5181   .dest_components = 1,
5182   .dest_bit_sizes = 0x20,
5183   .bit_size_src = -1,
5184   .num_indices = 0,
5185   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5186},
5187{
5188   .name = "load_sample_pos",
5189   .num_srcs = 0,
5190   .has_dest = true,
5191   .dest_components = 2,
5192   .dest_bit_sizes = 0x20,
5193   .bit_size_src = -1,
5194   .num_indices = 0,
5195   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5196},
5197{
5198   .name = "load_sample_pos_from_id",
5199   .num_srcs = 1,
5200   .src_components = {
5201      1
5202   },
5203   .has_dest = true,
5204   .dest_components = 2,
5205   .dest_bit_sizes = 0x0,
5206   .bit_size_src = -1,
5207   .num_indices = 0,
5208   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5209},
5210{
5211   .name = "load_sample_positions_pan",
5212   .num_srcs = 0,
5213   .has_dest = true,
5214   .dest_components = 1,
5215   .dest_bit_sizes = 0x40,
5216   .bit_size_src = -1,
5217   .num_indices = 0,
5218   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5219},
5220{
5221   .name = "load_sampler_lod_parameters_pan",
5222   .num_srcs = 1,
5223   .src_components = {
5224      1
5225   },
5226   .has_dest = true,
5227   .dest_components = 0,
5228   .dest_bit_sizes = 0x0,
5229   .bit_size_src = -1,
5230   .num_indices = 0,
5231   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5232},
5233{
5234   .name = "load_sbt_amd",
5235   .num_srcs = 0,
5236   .has_dest = true,
5237   .dest_components = 4,
5238   .dest_bit_sizes = 0x20,
5239   .bit_size_src = -1,
5240   .num_indices = 1,
5241   .indices = {
5242      NIR_INTRINSIC_BINDING,
5243   },
5244   .index_map = {
5245      [NIR_INTRINSIC_BINDING] = 1,
5246    },
5247   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5248},
5249{
5250   .name = "load_scratch",
5251   .num_srcs = 1,
5252   .src_components = {
5253      1
5254   },
5255   .has_dest = true,
5256   .dest_components = 0,
5257   .dest_bit_sizes = 0x0,
5258   .bit_size_src = -1,
5259   .num_indices = 2,
5260   .indices = {
5261      NIR_INTRINSIC_ALIGN_MUL,
5262      NIR_INTRINSIC_ALIGN_OFFSET,
5263   },
5264   .index_map = {
5265      [NIR_INTRINSIC_ALIGN_MUL] = 1,
5266      [NIR_INTRINSIC_ALIGN_OFFSET] = 2,
5267    },
5268   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
5269},
5270{
5271   .name = "load_scratch_base_ptr",
5272   .num_srcs = 0,
5273   .has_dest = true,
5274   .dest_components = 0,
5275   .dest_bit_sizes = 0x60,
5276   .bit_size_src = -1,
5277   .num_indices = 1,
5278   .indices = {
5279      NIR_INTRINSIC_BASE,
5280   },
5281   .index_map = {
5282      [NIR_INTRINSIC_BASE] = 1,
5283    },
5284   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5285},
5286{
5287   .name = "load_scratch_dxil",
5288   .num_srcs = 1,
5289   .src_components = {
5290      1
5291   },
5292   .has_dest = true,
5293   .dest_components = 0,
5294   .dest_bit_sizes = 0x0,
5295   .bit_size_src = -1,
5296   .num_indices = 0,
5297   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
5298},
5299{
5300   .name = "load_shader_query_enabled_amd",
5301   .num_srcs = 0,
5302   .has_dest = true,
5303   .dest_components = 1,
5304   .dest_bit_sizes = 0x1,
5305   .bit_size_src = -1,
5306   .num_indices = 0,
5307   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5308},
5309{
5310   .name = "load_shader_record_ptr",
5311   .num_srcs = 0,
5312   .has_dest = true,
5313   .dest_components = 1,
5314   .dest_bit_sizes = 0x40,
5315   .bit_size_src = -1,
5316   .num_indices = 0,
5317   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5318},
5319{
5320   .name = "load_shared",
5321   .num_srcs = 1,
5322   .src_components = {
5323      1
5324   },
5325   .has_dest = true,
5326   .dest_components = 0,
5327   .dest_bit_sizes = 0x0,
5328   .bit_size_src = -1,
5329   .num_indices = 3,
5330   .indices = {
5331      NIR_INTRINSIC_BASE,
5332      NIR_INTRINSIC_ALIGN_MUL,
5333      NIR_INTRINSIC_ALIGN_OFFSET,
5334   },
5335   .index_map = {
5336      [NIR_INTRINSIC_BASE] = 1,
5337      [NIR_INTRINSIC_ALIGN_MUL] = 2,
5338      [NIR_INTRINSIC_ALIGN_OFFSET] = 3,
5339    },
5340   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
5341},
5342{
5343   .name = "load_shared_base_ptr",
5344   .num_srcs = 0,
5345   .has_dest = true,
5346   .dest_components = 0,
5347   .dest_bit_sizes = 0x60,
5348   .bit_size_src = -1,
5349   .num_indices = 0,
5350   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5351},
5352{
5353   .name = "load_shared_block_intel",
5354   .num_srcs = 1,
5355   .src_components = {
5356      1
5357   },
5358   .has_dest = true,
5359   .dest_components = 0,
5360   .dest_bit_sizes = 0x0,
5361   .bit_size_src = -1,
5362   .num_indices = 3,
5363   .indices = {
5364      NIR_INTRINSIC_BASE,
5365      NIR_INTRINSIC_ALIGN_MUL,
5366      NIR_INTRINSIC_ALIGN_OFFSET,
5367   },
5368   .index_map = {
5369      [NIR_INTRINSIC_BASE] = 1,
5370      [NIR_INTRINSIC_ALIGN_MUL] = 2,
5371      [NIR_INTRINSIC_ALIGN_OFFSET] = 3,
5372    },
5373   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
5374},
5375{
5376   .name = "load_shared_dxil",
5377   .num_srcs = 1,
5378   .src_components = {
5379      1
5380   },
5381   .has_dest = true,
5382   .dest_components = 0,
5383   .dest_bit_sizes = 0x0,
5384   .bit_size_src = -1,
5385   .num_indices = 0,
5386   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
5387},
5388{
5389   .name = "load_shared_ir3",
5390   .num_srcs = 1,
5391   .src_components = {
5392      1
5393   },
5394   .has_dest = true,
5395   .dest_components = 0,
5396   .dest_bit_sizes = 0x0,
5397   .bit_size_src = -1,
5398   .num_indices = 3,
5399   .indices = {
5400      NIR_INTRINSIC_BASE,
5401      NIR_INTRINSIC_ALIGN_MUL,
5402      NIR_INTRINSIC_ALIGN_OFFSET,
5403   },
5404   .index_map = {
5405      [NIR_INTRINSIC_BASE] = 1,
5406      [NIR_INTRINSIC_ALIGN_MUL] = 2,
5407      [NIR_INTRINSIC_ALIGN_OFFSET] = 3,
5408    },
5409   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
5410},
5411{
5412   .name = "load_simd_width_intel",
5413   .num_srcs = 0,
5414   .has_dest = true,
5415   .dest_components = 1,
5416   .dest_bit_sizes = 0x20,
5417   .bit_size_src = -1,
5418   .num_indices = 0,
5419   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5420},
5421{
5422   .name = "load_size_ir3",
5423   .num_srcs = 0,
5424   .has_dest = true,
5425   .dest_components = 1,
5426   .dest_bit_sizes = 0x0,
5427   .bit_size_src = -1,
5428   .num_indices = 0,
5429   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5430},
5431{
5432   .name = "load_ssbo",
5433   .num_srcs = 2,
5434   .src_components = {
5435      -1, 1
5436   },
5437   .has_dest = true,
5438   .dest_components = 0,
5439   .dest_bit_sizes = 0x0,
5440   .bit_size_src = -1,
5441   .num_indices = 3,
5442   .indices = {
5443      NIR_INTRINSIC_ACCESS,
5444      NIR_INTRINSIC_ALIGN_MUL,
5445      NIR_INTRINSIC_ALIGN_OFFSET,
5446   },
5447   .index_map = {
5448      [NIR_INTRINSIC_ACCESS] = 1,
5449      [NIR_INTRINSIC_ALIGN_MUL] = 2,
5450      [NIR_INTRINSIC_ALIGN_OFFSET] = 3,
5451    },
5452   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
5453},
5454{
5455   .name = "load_ssbo_address",
5456   .num_srcs = 1,
5457   .src_components = {
5458      1
5459   },
5460   .has_dest = true,
5461   .dest_components = 0,
5462   .dest_bit_sizes = 0x0,
5463   .bit_size_src = -1,
5464   .num_indices = 0,
5465   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5466},
5467{
5468   .name = "load_ssbo_block_intel",
5469   .num_srcs = 2,
5470   .src_components = {
5471      -1, 1
5472   },
5473   .has_dest = true,
5474   .dest_components = 0,
5475   .dest_bit_sizes = 0x0,
5476   .bit_size_src = -1,
5477   .num_indices = 3,
5478   .indices = {
5479      NIR_INTRINSIC_ACCESS,
5480      NIR_INTRINSIC_ALIGN_MUL,
5481      NIR_INTRINSIC_ALIGN_OFFSET,
5482   },
5483   .index_map = {
5484      [NIR_INTRINSIC_ACCESS] = 1,
5485      [NIR_INTRINSIC_ALIGN_MUL] = 2,
5486      [NIR_INTRINSIC_ALIGN_OFFSET] = 3,
5487    },
5488   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
5489},
5490{
5491   .name = "load_ssbo_ir3",
5492   .num_srcs = 3,
5493   .src_components = {
5494      1, 1, 1
5495   },
5496   .has_dest = true,
5497   .dest_components = 0,
5498   .dest_bit_sizes = 0x0,
5499   .bit_size_src = -1,
5500   .num_indices = 3,
5501   .indices = {
5502      NIR_INTRINSIC_ACCESS,
5503      NIR_INTRINSIC_ALIGN_MUL,
5504      NIR_INTRINSIC_ALIGN_OFFSET,
5505   },
5506   .index_map = {
5507      [NIR_INTRINSIC_ACCESS] = 1,
5508      [NIR_INTRINSIC_ALIGN_MUL] = 2,
5509      [NIR_INTRINSIC_ALIGN_OFFSET] = 3,
5510    },
5511   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
5512},
5513{
5514   .name = "load_subgroup_eq_mask",
5515   .num_srcs = 0,
5516   .has_dest = true,
5517   .dest_components = 0,
5518   .dest_bit_sizes = 0x60,
5519   .bit_size_src = -1,
5520   .num_indices = 0,
5521   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5522},
5523{
5524   .name = "load_subgroup_ge_mask",
5525   .num_srcs = 0,
5526   .has_dest = true,
5527   .dest_components = 0,
5528   .dest_bit_sizes = 0x60,
5529   .bit_size_src = -1,
5530   .num_indices = 0,
5531   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5532},
5533{
5534   .name = "load_subgroup_gt_mask",
5535   .num_srcs = 0,
5536   .has_dest = true,
5537   .dest_components = 0,
5538   .dest_bit_sizes = 0x60,
5539   .bit_size_src = -1,
5540   .num_indices = 0,
5541   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5542},
5543{
5544   .name = "load_subgroup_id",
5545   .num_srcs = 0,
5546   .has_dest = true,
5547   .dest_components = 1,
5548   .dest_bit_sizes = 0x20,
5549   .bit_size_src = -1,
5550   .num_indices = 0,
5551   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5552},
5553{
5554   .name = "load_subgroup_id_shift_ir3",
5555   .num_srcs = 0,
5556   .has_dest = true,
5557   .dest_components = 1,
5558   .dest_bit_sizes = 0x20,
5559   .bit_size_src = -1,
5560   .num_indices = 0,
5561   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5562},
5563{
5564   .name = "load_subgroup_invocation",
5565   .num_srcs = 0,
5566   .has_dest = true,
5567   .dest_components = 1,
5568   .dest_bit_sizes = 0x20,
5569   .bit_size_src = -1,
5570   .num_indices = 0,
5571   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5572},
5573{
5574   .name = "load_subgroup_le_mask",
5575   .num_srcs = 0,
5576   .has_dest = true,
5577   .dest_components = 0,
5578   .dest_bit_sizes = 0x60,
5579   .bit_size_src = -1,
5580   .num_indices = 0,
5581   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5582},
5583{
5584   .name = "load_subgroup_lt_mask",
5585   .num_srcs = 0,
5586   .has_dest = true,
5587   .dest_components = 0,
5588   .dest_bit_sizes = 0x60,
5589   .bit_size_src = -1,
5590   .num_indices = 0,
5591   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5592},
5593{
5594   .name = "load_subgroup_size",
5595   .num_srcs = 0,
5596   .has_dest = true,
5597   .dest_components = 1,
5598   .dest_bit_sizes = 0x20,
5599   .bit_size_src = -1,
5600   .num_indices = 0,
5601   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5602},
5603{
5604   .name = "load_tcs_header_ir3",
5605   .num_srcs = 0,
5606   .has_dest = true,
5607   .dest_components = 1,
5608   .dest_bit_sizes = 0x20,
5609   .bit_size_src = -1,
5610   .num_indices = 0,
5611   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5612},
5613{
5614   .name = "load_tcs_in_param_base_r600",
5615   .num_srcs = 0,
5616   .has_dest = true,
5617   .dest_components = 4,
5618   .dest_bit_sizes = 0x20,
5619   .bit_size_src = -1,
5620   .num_indices = 0,
5621   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5622},
5623{
5624   .name = "load_tcs_num_patches_amd",
5625   .num_srcs = 0,
5626   .has_dest = true,
5627   .dest_components = 1,
5628   .dest_bit_sizes = 0x20,
5629   .bit_size_src = -1,
5630   .num_indices = 0,
5631   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5632},
5633{
5634   .name = "load_tcs_out_param_base_r600",
5635   .num_srcs = 0,
5636   .has_dest = true,
5637   .dest_components = 4,
5638   .dest_bit_sizes = 0x20,
5639   .bit_size_src = -1,
5640   .num_indices = 0,
5641   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5642},
5643{
5644   .name = "load_tcs_rel_patch_id_r600",
5645   .num_srcs = 0,
5646   .has_dest = true,
5647   .dest_components = 1,
5648   .dest_bit_sizes = 0x20,
5649   .bit_size_src = -1,
5650   .num_indices = 0,
5651   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5652},
5653{
5654   .name = "load_tcs_tess_factor_base_r600",
5655   .num_srcs = 0,
5656   .has_dest = true,
5657   .dest_components = 1,
5658   .dest_bit_sizes = 0x20,
5659   .bit_size_src = -1,
5660   .num_indices = 0,
5661   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5662},
5663{
5664   .name = "load_tess_coord",
5665   .num_srcs = 0,
5666   .has_dest = true,
5667   .dest_components = 3,
5668   .dest_bit_sizes = 0x20,
5669   .bit_size_src = -1,
5670   .num_indices = 0,
5671   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5672},
5673{
5674   .name = "load_tess_coord_r600",
5675   .num_srcs = 0,
5676   .has_dest = true,
5677   .dest_components = 2,
5678   .dest_bit_sizes = 0x20,
5679   .bit_size_src = -1,
5680   .num_indices = 0,
5681   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5682},
5683{
5684   .name = "load_tess_factor_base_ir3",
5685   .num_srcs = 0,
5686   .has_dest = true,
5687   .dest_components = 2,
5688   .dest_bit_sizes = 0x20,
5689   .bit_size_src = -1,
5690   .num_indices = 0,
5691   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5692},
5693{
5694   .name = "load_tess_level_inner",
5695   .num_srcs = 0,
5696   .has_dest = true,
5697   .dest_components = 2,
5698   .dest_bit_sizes = 0x20,
5699   .bit_size_src = -1,
5700   .num_indices = 0,
5701   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5702},
5703{
5704   .name = "load_tess_level_inner_default",
5705   .num_srcs = 0,
5706   .has_dest = true,
5707   .dest_components = 2,
5708   .dest_bit_sizes = 0x20,
5709   .bit_size_src = -1,
5710   .num_indices = 0,
5711   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5712},
5713{
5714   .name = "load_tess_level_outer",
5715   .num_srcs = 0,
5716   .has_dest = true,
5717   .dest_components = 4,
5718   .dest_bit_sizes = 0x20,
5719   .bit_size_src = -1,
5720   .num_indices = 0,
5721   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5722},
5723{
5724   .name = "load_tess_level_outer_default",
5725   .num_srcs = 0,
5726   .has_dest = true,
5727   .dest_components = 4,
5728   .dest_bit_sizes = 0x20,
5729   .bit_size_src = -1,
5730   .num_indices = 0,
5731   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5732},
5733{
5734   .name = "load_tess_param_base_ir3",
5735   .num_srcs = 0,
5736   .has_dest = true,
5737   .dest_components = 2,
5738   .dest_bit_sizes = 0x20,
5739   .bit_size_src = -1,
5740   .num_indices = 0,
5741   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5742},
5743{
5744   .name = "load_tess_rel_patch_id_amd",
5745   .num_srcs = 0,
5746   .has_dest = true,
5747   .dest_components = 1,
5748   .dest_bit_sizes = 0x20,
5749   .bit_size_src = -1,
5750   .num_indices = 0,
5751   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5752},
5753{
5754   .name = "load_texture_rect_scaling",
5755   .num_srcs = 1,
5756   .src_components = {
5757      1
5758   },
5759   .has_dest = true,
5760   .dest_components = 2,
5761   .dest_bit_sizes = 0x0,
5762   .bit_size_src = -1,
5763   .num_indices = 0,
5764   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5765},
5766{
5767   .name = "load_tlb_color_v3d",
5768   .num_srcs = 1,
5769   .src_components = {
5770      1
5771   },
5772   .has_dest = true,
5773   .dest_components = 0,
5774   .dest_bit_sizes = 0x0,
5775   .bit_size_src = -1,
5776   .num_indices = 2,
5777   .indices = {
5778      NIR_INTRINSIC_BASE,
5779      NIR_INTRINSIC_COMPONENT,
5780   },
5781   .index_map = {
5782      [NIR_INTRINSIC_BASE] = 1,
5783      [NIR_INTRINSIC_COMPONENT] = 2,
5784    },
5785   .flags = 0,
5786},
5787{
5788   .name = "load_ubo",
5789   .num_srcs = 2,
5790   .src_components = {
5791      -1, 1
5792   },
5793   .has_dest = true,
5794   .dest_components = 0,
5795   .dest_bit_sizes = 0x0,
5796   .bit_size_src = -1,
5797   .num_indices = 5,
5798   .indices = {
5799      NIR_INTRINSIC_ACCESS,
5800      NIR_INTRINSIC_ALIGN_MUL,
5801      NIR_INTRINSIC_ALIGN_OFFSET,
5802      NIR_INTRINSIC_RANGE_BASE,
5803      NIR_INTRINSIC_RANGE,
5804   },
5805   .index_map = {
5806      [NIR_INTRINSIC_ACCESS] = 1,
5807      [NIR_INTRINSIC_ALIGN_MUL] = 2,
5808      [NIR_INTRINSIC_ALIGN_OFFSET] = 3,
5809      [NIR_INTRINSIC_RANGE_BASE] = 4,
5810      [NIR_INTRINSIC_RANGE] = 5,
5811    },
5812   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5813},
5814{
5815   .name = "load_ubo_dxil",
5816   .num_srcs = 2,
5817   .src_components = {
5818      1, 1
5819   },
5820   .has_dest = true,
5821   .dest_components = 0,
5822   .dest_bit_sizes = 0x0,
5823   .bit_size_src = -1,
5824   .num_indices = 0,
5825   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5826},
5827{
5828   .name = "load_ubo_vec4",
5829   .num_srcs = 2,
5830   .src_components = {
5831      -1, 1
5832   },
5833   .has_dest = true,
5834   .dest_components = 0,
5835   .dest_bit_sizes = 0x0,
5836   .bit_size_src = -1,
5837   .num_indices = 2,
5838   .indices = {
5839      NIR_INTRINSIC_ACCESS,
5840      NIR_INTRINSIC_COMPONENT,
5841   },
5842   .index_map = {
5843      [NIR_INTRINSIC_ACCESS] = 1,
5844      [NIR_INTRINSIC_COMPONENT] = 2,
5845    },
5846   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5847},
5848{
5849   .name = "load_uniform",
5850   .num_srcs = 1,
5851   .src_components = {
5852      1
5853   },
5854   .has_dest = true,
5855   .dest_components = 0,
5856   .dest_bit_sizes = 0x0,
5857   .bit_size_src = -1,
5858   .num_indices = 3,
5859   .indices = {
5860      NIR_INTRINSIC_BASE,
5861      NIR_INTRINSIC_RANGE,
5862      NIR_INTRINSIC_DEST_TYPE,
5863   },
5864   .index_map = {
5865      [NIR_INTRINSIC_BASE] = 1,
5866      [NIR_INTRINSIC_RANGE] = 2,
5867      [NIR_INTRINSIC_DEST_TYPE] = 3,
5868    },
5869   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5870},
5871{
5872   .name = "load_user_clip_plane",
5873   .num_srcs = 0,
5874   .has_dest = true,
5875   .dest_components = 4,
5876   .dest_bit_sizes = 0x20,
5877   .bit_size_src = -1,
5878   .num_indices = 1,
5879   .indices = {
5880      NIR_INTRINSIC_UCP_ID,
5881   },
5882   .index_map = {
5883      [NIR_INTRINSIC_UCP_ID] = 1,
5884    },
5885   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5886},
5887{
5888   .name = "load_user_data_amd",
5889   .num_srcs = 0,
5890   .has_dest = true,
5891   .dest_components = 4,
5892   .dest_bit_sizes = 0x20,
5893   .bit_size_src = -1,
5894   .num_indices = 0,
5895   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5896},
5897{
5898   .name = "load_vertex_id",
5899   .num_srcs = 0,
5900   .has_dest = true,
5901   .dest_components = 1,
5902   .dest_bit_sizes = 0x20,
5903   .bit_size_src = -1,
5904   .num_indices = 0,
5905   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5906},
5907{
5908   .name = "load_vertex_id_zero_base",
5909   .num_srcs = 0,
5910   .has_dest = true,
5911   .dest_components = 1,
5912   .dest_bit_sizes = 0x20,
5913   .bit_size_src = -1,
5914   .num_indices = 0,
5915   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5916},
5917{
5918   .name = "load_view_index",
5919   .num_srcs = 0,
5920   .has_dest = true,
5921   .dest_components = 1,
5922   .dest_bit_sizes = 0x20,
5923   .bit_size_src = -1,
5924   .num_indices = 0,
5925   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5926},
5927{
5928   .name = "load_viewport_offset",
5929   .num_srcs = 0,
5930   .has_dest = true,
5931   .dest_components = 3,
5932   .dest_bit_sizes = 0x20,
5933   .bit_size_src = -1,
5934   .num_indices = 0,
5935   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5936},
5937{
5938   .name = "load_viewport_scale",
5939   .num_srcs = 0,
5940   .has_dest = true,
5941   .dest_components = 3,
5942   .dest_bit_sizes = 0x20,
5943   .bit_size_src = -1,
5944   .num_indices = 0,
5945   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5946},
5947{
5948   .name = "load_viewport_x_offset",
5949   .num_srcs = 0,
5950   .has_dest = true,
5951   .dest_components = 1,
5952   .dest_bit_sizes = 0x20,
5953   .bit_size_src = -1,
5954   .num_indices = 0,
5955   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5956},
5957{
5958   .name = "load_viewport_x_scale",
5959   .num_srcs = 0,
5960   .has_dest = true,
5961   .dest_components = 1,
5962   .dest_bit_sizes = 0x20,
5963   .bit_size_src = -1,
5964   .num_indices = 0,
5965   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5966},
5967{
5968   .name = "load_viewport_y_offset",
5969   .num_srcs = 0,
5970   .has_dest = true,
5971   .dest_components = 1,
5972   .dest_bit_sizes = 0x20,
5973   .bit_size_src = -1,
5974   .num_indices = 0,
5975   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5976},
5977{
5978   .name = "load_viewport_y_scale",
5979   .num_srcs = 0,
5980   .has_dest = true,
5981   .dest_components = 1,
5982   .dest_bit_sizes = 0x20,
5983   .bit_size_src = -1,
5984   .num_indices = 0,
5985   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5986},
5987{
5988   .name = "load_viewport_z_offset",
5989   .num_srcs = 0,
5990   .has_dest = true,
5991   .dest_components = 1,
5992   .dest_bit_sizes = 0x20,
5993   .bit_size_src = -1,
5994   .num_indices = 0,
5995   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
5996},
5997{
5998   .name = "load_viewport_z_scale",
5999   .num_srcs = 0,
6000   .has_dest = true,
6001   .dest_components = 1,
6002   .dest_bit_sizes = 0x20,
6003   .bit_size_src = -1,
6004   .num_indices = 0,
6005   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
6006},
6007{
6008   .name = "load_vs_primitive_stride_ir3",
6009   .num_srcs = 0,
6010   .has_dest = true,
6011   .dest_components = 1,
6012   .dest_bit_sizes = 0x20,
6013   .bit_size_src = -1,
6014   .num_indices = 0,
6015   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
6016},
6017{
6018   .name = "load_vs_vertex_stride_ir3",
6019   .num_srcs = 0,
6020   .has_dest = true,
6021   .dest_components = 1,
6022   .dest_bit_sizes = 0x20,
6023   .bit_size_src = -1,
6024   .num_indices = 0,
6025   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
6026},
6027{
6028   .name = "load_vulkan_descriptor",
6029   .num_srcs = 1,
6030   .src_components = {
6031      -1
6032   },
6033   .has_dest = true,
6034   .dest_components = 0,
6035   .dest_bit_sizes = 0x0,
6036   .bit_size_src = -1,
6037   .num_indices = 1,
6038   .indices = {
6039      NIR_INTRINSIC_DESC_TYPE,
6040   },
6041   .index_map = {
6042      [NIR_INTRINSIC_DESC_TYPE] = 1,
6043    },
6044   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
6045},
6046{
6047   .name = "load_work_dim",
6048   .num_srcs = 0,
6049   .has_dest = true,
6050   .dest_components = 1,
6051   .dest_bit_sizes = 0x20,
6052   .bit_size_src = -1,
6053   .num_indices = 0,
6054   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
6055},
6056{
6057   .name = "load_workgroup_id",
6058   .num_srcs = 0,
6059   .has_dest = true,
6060   .dest_components = 3,
6061   .dest_bit_sizes = 0x60,
6062   .bit_size_src = -1,
6063   .num_indices = 0,
6064   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
6065},
6066{
6067   .name = "load_workgroup_id_zero_base",
6068   .num_srcs = 0,
6069   .has_dest = true,
6070   .dest_components = 3,
6071   .dest_bit_sizes = 0x20,
6072   .bit_size_src = -1,
6073   .num_indices = 0,
6074   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
6075},
6076{
6077   .name = "load_workgroup_num_input_primitives_amd",
6078   .num_srcs = 0,
6079   .has_dest = true,
6080   .dest_components = 1,
6081   .dest_bit_sizes = 0x20,
6082   .bit_size_src = -1,
6083   .num_indices = 0,
6084   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
6085},
6086{
6087   .name = "load_workgroup_num_input_vertices_amd",
6088   .num_srcs = 0,
6089   .has_dest = true,
6090   .dest_components = 1,
6091   .dest_bit_sizes = 0x20,
6092   .bit_size_src = -1,
6093   .num_indices = 0,
6094   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
6095},
6096{
6097   .name = "load_workgroup_size",
6098   .num_srcs = 0,
6099   .has_dest = true,
6100   .dest_components = 3,
6101   .dest_bit_sizes = 0x20,
6102   .bit_size_src = -1,
6103   .num_indices = 0,
6104   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
6105},
6106{
6107   .name = "masked_swizzle_amd",
6108   .num_srcs = 1,
6109   .src_components = {
6110      0
6111   },
6112   .has_dest = true,
6113   .dest_components = 0,
6114   .dest_bit_sizes = 0x0,
6115   .bit_size_src = 0,
6116   .num_indices = 1,
6117   .indices = {
6118      NIR_INTRINSIC_SWIZZLE_MASK,
6119   },
6120   .index_map = {
6121      [NIR_INTRINSIC_SWIZZLE_MASK] = 1,
6122    },
6123   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
6124},
6125{
6126   .name = "mbcnt_amd",
6127   .num_srcs = 2,
6128   .src_components = {
6129      1, 1
6130   },
6131   .has_dest = true,
6132   .dest_components = 1,
6133   .dest_bit_sizes = 0x20,
6134   .bit_size_src = -1,
6135   .num_indices = 0,
6136   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
6137},
6138{
6139   .name = "memcpy_deref",
6140   .num_srcs = 3,
6141   .src_components = {
6142      -1, -1, 1
6143   },
6144   .has_dest = false,
6145   .dest_components = 0,
6146   .dest_bit_sizes = 0x0,
6147   .bit_size_src = -1,
6148   .num_indices = 2,
6149   .indices = {
6150      NIR_INTRINSIC_DST_ACCESS,
6151      NIR_INTRINSIC_SRC_ACCESS,
6152   },
6153   .index_map = {
6154      [NIR_INTRINSIC_DST_ACCESS] = 1,
6155      [NIR_INTRINSIC_SRC_ACCESS] = 2,
6156    },
6157   .flags = 0,
6158},
6159{
6160   .name = "memory_barrier",
6161   .num_srcs = 0,
6162   .has_dest = false,
6163   .dest_components = 0,
6164   .dest_bit_sizes = 0x0,
6165   .bit_size_src = -1,
6166   .num_indices = 0,
6167   .flags = 0,
6168},
6169{
6170   .name = "memory_barrier_atomic_counter",
6171   .num_srcs = 0,
6172   .has_dest = false,
6173   .dest_components = 0,
6174   .dest_bit_sizes = 0x0,
6175   .bit_size_src = -1,
6176   .num_indices = 0,
6177   .flags = 0,
6178},
6179{
6180   .name = "memory_barrier_buffer",
6181   .num_srcs = 0,
6182   .has_dest = false,
6183   .dest_components = 0,
6184   .dest_bit_sizes = 0x0,
6185   .bit_size_src = -1,
6186   .num_indices = 0,
6187   .flags = 0,
6188},
6189{
6190   .name = "memory_barrier_image",
6191   .num_srcs = 0,
6192   .has_dest = false,
6193   .dest_components = 0,
6194   .dest_bit_sizes = 0x0,
6195   .bit_size_src = -1,
6196   .num_indices = 0,
6197   .flags = 0,
6198},
6199{
6200   .name = "memory_barrier_shared",
6201   .num_srcs = 0,
6202   .has_dest = false,
6203   .dest_components = 0,
6204   .dest_bit_sizes = 0x0,
6205   .bit_size_src = -1,
6206   .num_indices = 0,
6207   .flags = 0,
6208},
6209{
6210   .name = "memory_barrier_tcs_patch",
6211   .num_srcs = 0,
6212   .has_dest = false,
6213   .dest_components = 0,
6214   .dest_bit_sizes = 0x0,
6215   .bit_size_src = -1,
6216   .num_indices = 0,
6217   .flags = 0,
6218},
6219{
6220   .name = "nop",
6221   .num_srcs = 0,
6222   .has_dest = false,
6223   .dest_components = 0,
6224   .dest_bit_sizes = 0x0,
6225   .bit_size_src = -1,
6226   .num_indices = 0,
6227   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
6228},
6229{
6230   .name = "overwrite_tes_arguments_amd",
6231   .num_srcs = 4,
6232   .src_components = {
6233      1, 1, 1, 1
6234   },
6235   .has_dest = false,
6236   .dest_components = 0,
6237   .dest_bit_sizes = 0x0,
6238   .bit_size_src = -1,
6239   .num_indices = 0,
6240   .flags = 0,
6241},
6242{
6243   .name = "overwrite_vs_arguments_amd",
6244   .num_srcs = 2,
6245   .src_components = {
6246      1, 1
6247   },
6248   .has_dest = false,
6249   .dest_components = 0,
6250   .dest_bit_sizes = 0x0,
6251   .bit_size_src = -1,
6252   .num_indices = 0,
6253   .flags = 0,
6254},
6255{
6256   .name = "printf",
6257   .num_srcs = 2,
6258   .src_components = {
6259      1, 1
6260   },
6261   .has_dest = true,
6262   .dest_components = 1,
6263   .dest_bit_sizes = 0x20,
6264   .bit_size_src = -1,
6265   .num_indices = 0,
6266   .flags = 0,
6267},
6268{
6269   .name = "quad_broadcast",
6270   .num_srcs = 2,
6271   .src_components = {
6272      0, 1
6273   },
6274   .has_dest = true,
6275   .dest_components = 0,
6276   .dest_bit_sizes = 0x0,
6277   .bit_size_src = -1,
6278   .num_indices = 0,
6279   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
6280},
6281{
6282   .name = "quad_swap_diagonal",
6283   .num_srcs = 1,
6284   .src_components = {
6285      0
6286   },
6287   .has_dest = true,
6288   .dest_components = 0,
6289   .dest_bit_sizes = 0x0,
6290   .bit_size_src = -1,
6291   .num_indices = 0,
6292   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
6293},
6294{
6295   .name = "quad_swap_horizontal",
6296   .num_srcs = 1,
6297   .src_components = {
6298      0
6299   },
6300   .has_dest = true,
6301   .dest_components = 0,
6302   .dest_bit_sizes = 0x0,
6303   .bit_size_src = -1,
6304   .num_indices = 0,
6305   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
6306},
6307{
6308   .name = "quad_swap_vertical",
6309   .num_srcs = 1,
6310   .src_components = {
6311      0
6312   },
6313   .has_dest = true,
6314   .dest_components = 0,
6315   .dest_bit_sizes = 0x0,
6316   .bit_size_src = -1,
6317   .num_indices = 0,
6318   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
6319},
6320{
6321   .name = "quad_swizzle_amd",
6322   .num_srcs = 1,
6323   .src_components = {
6324      0
6325   },
6326   .has_dest = true,
6327   .dest_components = 0,
6328   .dest_bit_sizes = 0x0,
6329   .bit_size_src = 0,
6330   .num_indices = 1,
6331   .indices = {
6332      NIR_INTRINSIC_SWIZZLE_MASK,
6333   },
6334   .index_map = {
6335      [NIR_INTRINSIC_SWIZZLE_MASK] = 1,
6336    },
6337   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
6338},
6339{
6340   .name = "read_first_invocation",
6341   .num_srcs = 1,
6342   .src_components = {
6343      0
6344   },
6345   .has_dest = true,
6346   .dest_components = 0,
6347   .dest_bit_sizes = 0x0,
6348   .bit_size_src = 0,
6349   .num_indices = 0,
6350   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
6351},
6352{
6353   .name = "read_invocation",
6354   .num_srcs = 2,
6355   .src_components = {
6356      0, 1
6357   },
6358   .has_dest = true,
6359   .dest_components = 0,
6360   .dest_bit_sizes = 0x0,
6361   .bit_size_src = 0,
6362   .num_indices = 0,
6363   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
6364},
6365{
6366   .name = "read_invocation_cond_ir3",
6367   .num_srcs = 2,
6368   .src_components = {
6369      0, 1
6370   },
6371   .has_dest = true,
6372   .dest_components = 0,
6373   .dest_bit_sizes = 0x0,
6374   .bit_size_src = -1,
6375   .num_indices = 0,
6376   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
6377},
6378{
6379   .name = "reduce",
6380   .num_srcs = 1,
6381   .src_components = {
6382      0
6383   },
6384   .has_dest = true,
6385   .dest_components = 0,
6386   .dest_bit_sizes = 0x0,
6387   .bit_size_src = 0,
6388   .num_indices = 2,
6389   .indices = {
6390      NIR_INTRINSIC_REDUCTION_OP,
6391      NIR_INTRINSIC_CLUSTER_SIZE,
6392   },
6393   .index_map = {
6394      [NIR_INTRINSIC_REDUCTION_OP] = 1,
6395      [NIR_INTRINSIC_CLUSTER_SIZE] = 2,
6396    },
6397   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
6398},
6399{
6400   .name = "report_ray_intersection",
6401   .num_srcs = 2,
6402   .src_components = {
6403      1, 1
6404   },
6405   .has_dest = true,
6406   .dest_components = 1,
6407   .dest_bit_sizes = 0x0,
6408   .bit_size_src = -1,
6409   .num_indices = 0,
6410   .flags = 0,
6411},
6412{
6413   .name = "rt_execute_callable",
6414   .num_srcs = 2,
6415   .src_components = {
6416      1, -1
6417   },
6418   .has_dest = false,
6419   .dest_components = 0,
6420   .dest_bit_sizes = 0x0,
6421   .bit_size_src = -1,
6422   .num_indices = 2,
6423   .indices = {
6424      NIR_INTRINSIC_CALL_IDX,
6425      NIR_INTRINSIC_STACK_SIZE,
6426   },
6427   .index_map = {
6428      [NIR_INTRINSIC_CALL_IDX] = 1,
6429      [NIR_INTRINSIC_STACK_SIZE] = 2,
6430    },
6431   .flags = 0,
6432},
6433{
6434   .name = "rt_resume",
6435   .num_srcs = 0,
6436   .has_dest = false,
6437   .dest_components = 0,
6438   .dest_bit_sizes = 0x0,
6439   .bit_size_src = -1,
6440   .num_indices = 2,
6441   .indices = {
6442      NIR_INTRINSIC_CALL_IDX,
6443      NIR_INTRINSIC_STACK_SIZE,
6444   },
6445   .index_map = {
6446      [NIR_INTRINSIC_CALL_IDX] = 1,
6447      [NIR_INTRINSIC_STACK_SIZE] = 2,
6448    },
6449   .flags = 0,
6450},
6451{
6452   .name = "rt_return_amd",
6453   .num_srcs = 0,
6454   .has_dest = false,
6455   .dest_components = 0,
6456   .dest_bit_sizes = 0x0,
6457   .bit_size_src = -1,
6458   .num_indices = 0,
6459   .flags = 0,
6460},
6461{
6462   .name = "rt_trace_ray",
6463   .num_srcs = 11,
6464   .src_components = {
6465      -1, 1, 1, 1, 1, 1, 3, 1, 3, 1, -1
6466   },
6467   .has_dest = false,
6468   .dest_components = 0,
6469   .dest_bit_sizes = 0x0,
6470   .bit_size_src = -1,
6471   .num_indices = 2,
6472   .indices = {
6473      NIR_INTRINSIC_CALL_IDX,
6474      NIR_INTRINSIC_STACK_SIZE,
6475   },
6476   .index_map = {
6477      [NIR_INTRINSIC_CALL_IDX] = 1,
6478      [NIR_INTRINSIC_STACK_SIZE] = 2,
6479    },
6480   .flags = 0,
6481},
6482{
6483   .name = "scoped_barrier",
6484   .num_srcs = 0,
6485   .has_dest = false,
6486   .dest_components = 0,
6487   .dest_bit_sizes = 0x0,
6488   .bit_size_src = -1,
6489   .num_indices = 4,
6490   .indices = {
6491      NIR_INTRINSIC_EXECUTION_SCOPE,
6492      NIR_INTRINSIC_MEMORY_SCOPE,
6493      NIR_INTRINSIC_MEMORY_SEMANTICS,
6494      NIR_INTRINSIC_MEMORY_MODES,
6495   },
6496   .index_map = {
6497      [NIR_INTRINSIC_EXECUTION_SCOPE] = 1,
6498      [NIR_INTRINSIC_MEMORY_SCOPE] = 2,
6499      [NIR_INTRINSIC_MEMORY_SEMANTICS] = 3,
6500      [NIR_INTRINSIC_MEMORY_MODES] = 4,
6501    },
6502   .flags = 0,
6503},
6504{
6505   .name = "set_vertex_and_primitive_count",
6506   .num_srcs = 2,
6507   .src_components = {
6508      1, 1
6509   },
6510   .has_dest = false,
6511   .dest_components = 0,
6512   .dest_bit_sizes = 0x0,
6513   .bit_size_src = -1,
6514   .num_indices = 1,
6515   .indices = {
6516      NIR_INTRINSIC_STREAM_ID,
6517   },
6518   .index_map = {
6519      [NIR_INTRINSIC_STREAM_ID] = 1,
6520    },
6521   .flags = 0,
6522},
6523{
6524   .name = "shader_clock",
6525   .num_srcs = 0,
6526   .has_dest = true,
6527   .dest_components = 2,
6528   .dest_bit_sizes = 0x20,
6529   .bit_size_src = -1,
6530   .num_indices = 1,
6531   .indices = {
6532      NIR_INTRINSIC_MEMORY_SCOPE,
6533   },
6534   .index_map = {
6535      [NIR_INTRINSIC_MEMORY_SCOPE] = 1,
6536    },
6537   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
6538},
6539{
6540   .name = "shared_atomic_add",
6541   .num_srcs = 2,
6542   .src_components = {
6543      1, 1
6544   },
6545   .has_dest = true,
6546   .dest_components = 1,
6547   .dest_bit_sizes = 0x0,
6548   .bit_size_src = -1,
6549   .num_indices = 1,
6550   .indices = {
6551      NIR_INTRINSIC_BASE,
6552   },
6553   .index_map = {
6554      [NIR_INTRINSIC_BASE] = 1,
6555    },
6556   .flags = 0,
6557},
6558{
6559   .name = "shared_atomic_add_dxil",
6560   .num_srcs = 2,
6561   .src_components = {
6562      1, 1
6563   },
6564   .has_dest = true,
6565   .dest_components = 1,
6566   .dest_bit_sizes = 0x0,
6567   .bit_size_src = -1,
6568   .num_indices = 0,
6569   .flags = 0,
6570},
6571{
6572   .name = "shared_atomic_and",
6573   .num_srcs = 2,
6574   .src_components = {
6575      1, 1
6576   },
6577   .has_dest = true,
6578   .dest_components = 1,
6579   .dest_bit_sizes = 0x0,
6580   .bit_size_src = -1,
6581   .num_indices = 1,
6582   .indices = {
6583      NIR_INTRINSIC_BASE,
6584   },
6585   .index_map = {
6586      [NIR_INTRINSIC_BASE] = 1,
6587    },
6588   .flags = 0,
6589},
6590{
6591   .name = "shared_atomic_and_dxil",
6592   .num_srcs = 2,
6593   .src_components = {
6594      1, 1
6595   },
6596   .has_dest = true,
6597   .dest_components = 1,
6598   .dest_bit_sizes = 0x0,
6599   .bit_size_src = -1,
6600   .num_indices = 0,
6601   .flags = 0,
6602},
6603{
6604   .name = "shared_atomic_comp_swap",
6605   .num_srcs = 3,
6606   .src_components = {
6607      1, 1, 1
6608   },
6609   .has_dest = true,
6610   .dest_components = 1,
6611   .dest_bit_sizes = 0x0,
6612   .bit_size_src = -1,
6613   .num_indices = 1,
6614   .indices = {
6615      NIR_INTRINSIC_BASE,
6616   },
6617   .index_map = {
6618      [NIR_INTRINSIC_BASE] = 1,
6619    },
6620   .flags = 0,
6621},
6622{
6623   .name = "shared_atomic_comp_swap_dxil",
6624   .num_srcs = 3,
6625   .src_components = {
6626      1, 1, 1
6627   },
6628   .has_dest = true,
6629   .dest_components = 1,
6630   .dest_bit_sizes = 0x0,
6631   .bit_size_src = -1,
6632   .num_indices = 0,
6633   .flags = 0,
6634},
6635{
6636   .name = "shared_atomic_exchange",
6637   .num_srcs = 2,
6638   .src_components = {
6639      1, 1
6640   },
6641   .has_dest = true,
6642   .dest_components = 1,
6643   .dest_bit_sizes = 0x0,
6644   .bit_size_src = -1,
6645   .num_indices = 1,
6646   .indices = {
6647      NIR_INTRINSIC_BASE,
6648   },
6649   .index_map = {
6650      [NIR_INTRINSIC_BASE] = 1,
6651    },
6652   .flags = 0,
6653},
6654{
6655   .name = "shared_atomic_exchange_dxil",
6656   .num_srcs = 2,
6657   .src_components = {
6658      1, 1
6659   },
6660   .has_dest = true,
6661   .dest_components = 1,
6662   .dest_bit_sizes = 0x0,
6663   .bit_size_src = -1,
6664   .num_indices = 0,
6665   .flags = 0,
6666},
6667{
6668   .name = "shared_atomic_fadd",
6669   .num_srcs = 2,
6670   .src_components = {
6671      1, 1
6672   },
6673   .has_dest = true,
6674   .dest_components = 1,
6675   .dest_bit_sizes = 0x0,
6676   .bit_size_src = -1,
6677   .num_indices = 1,
6678   .indices = {
6679      NIR_INTRINSIC_BASE,
6680   },
6681   .index_map = {
6682      [NIR_INTRINSIC_BASE] = 1,
6683    },
6684   .flags = 0,
6685},
6686{
6687   .name = "shared_atomic_fcomp_swap",
6688   .num_srcs = 3,
6689   .src_components = {
6690      1, 1, 1
6691   },
6692   .has_dest = true,
6693   .dest_components = 1,
6694   .dest_bit_sizes = 0x0,
6695   .bit_size_src = -1,
6696   .num_indices = 1,
6697   .indices = {
6698      NIR_INTRINSIC_BASE,
6699   },
6700   .index_map = {
6701      [NIR_INTRINSIC_BASE] = 1,
6702    },
6703   .flags = 0,
6704},
6705{
6706   .name = "shared_atomic_fmax",
6707   .num_srcs = 2,
6708   .src_components = {
6709      1, 1
6710   },
6711   .has_dest = true,
6712   .dest_components = 1,
6713   .dest_bit_sizes = 0x0,
6714   .bit_size_src = -1,
6715   .num_indices = 1,
6716   .indices = {
6717      NIR_INTRINSIC_BASE,
6718   },
6719   .index_map = {
6720      [NIR_INTRINSIC_BASE] = 1,
6721    },
6722   .flags = 0,
6723},
6724{
6725   .name = "shared_atomic_fmin",
6726   .num_srcs = 2,
6727   .src_components = {
6728      1, 1
6729   },
6730   .has_dest = true,
6731   .dest_components = 1,
6732   .dest_bit_sizes = 0x0,
6733   .bit_size_src = -1,
6734   .num_indices = 1,
6735   .indices = {
6736      NIR_INTRINSIC_BASE,
6737   },
6738   .index_map = {
6739      [NIR_INTRINSIC_BASE] = 1,
6740    },
6741   .flags = 0,
6742},
6743{
6744   .name = "shared_atomic_imax",
6745   .num_srcs = 2,
6746   .src_components = {
6747      1, 1
6748   },
6749   .has_dest = true,
6750   .dest_components = 1,
6751   .dest_bit_sizes = 0x0,
6752   .bit_size_src = -1,
6753   .num_indices = 1,
6754   .indices = {
6755      NIR_INTRINSIC_BASE,
6756   },
6757   .index_map = {
6758      [NIR_INTRINSIC_BASE] = 1,
6759    },
6760   .flags = 0,
6761},
6762{
6763   .name = "shared_atomic_imax_dxil",
6764   .num_srcs = 2,
6765   .src_components = {
6766      1, 1
6767   },
6768   .has_dest = true,
6769   .dest_components = 1,
6770   .dest_bit_sizes = 0x0,
6771   .bit_size_src = -1,
6772   .num_indices = 0,
6773   .flags = 0,
6774},
6775{
6776   .name = "shared_atomic_imin",
6777   .num_srcs = 2,
6778   .src_components = {
6779      1, 1
6780   },
6781   .has_dest = true,
6782   .dest_components = 1,
6783   .dest_bit_sizes = 0x0,
6784   .bit_size_src = -1,
6785   .num_indices = 1,
6786   .indices = {
6787      NIR_INTRINSIC_BASE,
6788   },
6789   .index_map = {
6790      [NIR_INTRINSIC_BASE] = 1,
6791    },
6792   .flags = 0,
6793},
6794{
6795   .name = "shared_atomic_imin_dxil",
6796   .num_srcs = 2,
6797   .src_components = {
6798      1, 1
6799   },
6800   .has_dest = true,
6801   .dest_components = 1,
6802   .dest_bit_sizes = 0x0,
6803   .bit_size_src = -1,
6804   .num_indices = 0,
6805   .flags = 0,
6806},
6807{
6808   .name = "shared_atomic_or",
6809   .num_srcs = 2,
6810   .src_components = {
6811      1, 1
6812   },
6813   .has_dest = true,
6814   .dest_components = 1,
6815   .dest_bit_sizes = 0x0,
6816   .bit_size_src = -1,
6817   .num_indices = 1,
6818   .indices = {
6819      NIR_INTRINSIC_BASE,
6820   },
6821   .index_map = {
6822      [NIR_INTRINSIC_BASE] = 1,
6823    },
6824   .flags = 0,
6825},
6826{
6827   .name = "shared_atomic_or_dxil",
6828   .num_srcs = 2,
6829   .src_components = {
6830      1, 1
6831   },
6832   .has_dest = true,
6833   .dest_components = 1,
6834   .dest_bit_sizes = 0x0,
6835   .bit_size_src = -1,
6836   .num_indices = 0,
6837   .flags = 0,
6838},
6839{
6840   .name = "shared_atomic_umax",
6841   .num_srcs = 2,
6842   .src_components = {
6843      1, 1
6844   },
6845   .has_dest = true,
6846   .dest_components = 1,
6847   .dest_bit_sizes = 0x0,
6848   .bit_size_src = -1,
6849   .num_indices = 1,
6850   .indices = {
6851      NIR_INTRINSIC_BASE,
6852   },
6853   .index_map = {
6854      [NIR_INTRINSIC_BASE] = 1,
6855    },
6856   .flags = 0,
6857},
6858{
6859   .name = "shared_atomic_umax_dxil",
6860   .num_srcs = 2,
6861   .src_components = {
6862      1, 1
6863   },
6864   .has_dest = true,
6865   .dest_components = 1,
6866   .dest_bit_sizes = 0x0,
6867   .bit_size_src = -1,
6868   .num_indices = 0,
6869   .flags = 0,
6870},
6871{
6872   .name = "shared_atomic_umin",
6873   .num_srcs = 2,
6874   .src_components = {
6875      1, 1
6876   },
6877   .has_dest = true,
6878   .dest_components = 1,
6879   .dest_bit_sizes = 0x0,
6880   .bit_size_src = -1,
6881   .num_indices = 1,
6882   .indices = {
6883      NIR_INTRINSIC_BASE,
6884   },
6885   .index_map = {
6886      [NIR_INTRINSIC_BASE] = 1,
6887    },
6888   .flags = 0,
6889},
6890{
6891   .name = "shared_atomic_umin_dxil",
6892   .num_srcs = 2,
6893   .src_components = {
6894      1, 1
6895   },
6896   .has_dest = true,
6897   .dest_components = 1,
6898   .dest_bit_sizes = 0x0,
6899   .bit_size_src = -1,
6900   .num_indices = 0,
6901   .flags = 0,
6902},
6903{
6904   .name = "shared_atomic_xor",
6905   .num_srcs = 2,
6906   .src_components = {
6907      1, 1
6908   },
6909   .has_dest = true,
6910   .dest_components = 1,
6911   .dest_bit_sizes = 0x0,
6912   .bit_size_src = -1,
6913   .num_indices = 1,
6914   .indices = {
6915      NIR_INTRINSIC_BASE,
6916   },
6917   .index_map = {
6918      [NIR_INTRINSIC_BASE] = 1,
6919    },
6920   .flags = 0,
6921},
6922{
6923   .name = "shared_atomic_xor_dxil",
6924   .num_srcs = 2,
6925   .src_components = {
6926      1, 1
6927   },
6928   .has_dest = true,
6929   .dest_components = 1,
6930   .dest_bit_sizes = 0x0,
6931   .bit_size_src = -1,
6932   .num_indices = 0,
6933   .flags = 0,
6934},
6935{
6936   .name = "shuffle",
6937   .num_srcs = 2,
6938   .src_components = {
6939      0, 1
6940   },
6941   .has_dest = true,
6942   .dest_components = 0,
6943   .dest_bit_sizes = 0x0,
6944   .bit_size_src = 0,
6945   .num_indices = 0,
6946   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
6947},
6948{
6949   .name = "shuffle_down",
6950   .num_srcs = 2,
6951   .src_components = {
6952      0, 1
6953   },
6954   .has_dest = true,
6955   .dest_components = 0,
6956   .dest_bit_sizes = 0x0,
6957   .bit_size_src = 0,
6958   .num_indices = 0,
6959   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
6960},
6961{
6962   .name = "shuffle_up",
6963   .num_srcs = 2,
6964   .src_components = {
6965      0, 1
6966   },
6967   .has_dest = true,
6968   .dest_components = 0,
6969   .dest_bit_sizes = 0x0,
6970   .bit_size_src = 0,
6971   .num_indices = 0,
6972   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
6973},
6974{
6975   .name = "shuffle_xor",
6976   .num_srcs = 2,
6977   .src_components = {
6978      0, 1
6979   },
6980   .has_dest = true,
6981   .dest_components = 0,
6982   .dest_bit_sizes = 0x0,
6983   .bit_size_src = 0,
6984   .num_indices = 0,
6985   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
6986},
6987{
6988   .name = "sparse_residency_code_and",
6989   .num_srcs = 2,
6990   .src_components = {
6991      1, 1
6992   },
6993   .has_dest = true,
6994   .dest_components = 1,
6995   .dest_bit_sizes = 0x20,
6996   .bit_size_src = -1,
6997   .num_indices = 0,
6998   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
6999},
7000{
7001   .name = "ssbo_atomic_add",
7002   .num_srcs = 3,
7003   .src_components = {
7004      -1, 1, 1
7005   },
7006   .has_dest = true,
7007   .dest_components = 1,
7008   .dest_bit_sizes = 0x0,
7009   .bit_size_src = -1,
7010   .num_indices = 1,
7011   .indices = {
7012      NIR_INTRINSIC_ACCESS,
7013   },
7014   .index_map = {
7015      [NIR_INTRINSIC_ACCESS] = 1,
7016    },
7017   .flags = 0,
7018},
7019{
7020   .name = "ssbo_atomic_add_ir3",
7021   .num_srcs = 4,
7022   .src_components = {
7023      1, 1, 1, 1
7024   },
7025   .has_dest = true,
7026   .dest_components = 1,
7027   .dest_bit_sizes = 0x0,
7028   .bit_size_src = -1,
7029   .num_indices = 1,
7030   .indices = {
7031      NIR_INTRINSIC_ACCESS,
7032   },
7033   .index_map = {
7034      [NIR_INTRINSIC_ACCESS] = 1,
7035    },
7036   .flags = 0,
7037},
7038{
7039   .name = "ssbo_atomic_and",
7040   .num_srcs = 3,
7041   .src_components = {
7042      -1, 1, 1
7043   },
7044   .has_dest = true,
7045   .dest_components = 1,
7046   .dest_bit_sizes = 0x0,
7047   .bit_size_src = -1,
7048   .num_indices = 1,
7049   .indices = {
7050      NIR_INTRINSIC_ACCESS,
7051   },
7052   .index_map = {
7053      [NIR_INTRINSIC_ACCESS] = 1,
7054    },
7055   .flags = 0,
7056},
7057{
7058   .name = "ssbo_atomic_and_ir3",
7059   .num_srcs = 4,
7060   .src_components = {
7061      1, 1, 1, 1
7062   },
7063   .has_dest = true,
7064   .dest_components = 1,
7065   .dest_bit_sizes = 0x0,
7066   .bit_size_src = -1,
7067   .num_indices = 1,
7068   .indices = {
7069      NIR_INTRINSIC_ACCESS,
7070   },
7071   .index_map = {
7072      [NIR_INTRINSIC_ACCESS] = 1,
7073    },
7074   .flags = 0,
7075},
7076{
7077   .name = "ssbo_atomic_comp_swap",
7078   .num_srcs = 4,
7079   .src_components = {
7080      -1, 1, 1, 1
7081   },
7082   .has_dest = true,
7083   .dest_components = 1,
7084   .dest_bit_sizes = 0x0,
7085   .bit_size_src = -1,
7086   .num_indices = 1,
7087   .indices = {
7088      NIR_INTRINSIC_ACCESS,
7089   },
7090   .index_map = {
7091      [NIR_INTRINSIC_ACCESS] = 1,
7092    },
7093   .flags = 0,
7094},
7095{
7096   .name = "ssbo_atomic_comp_swap_ir3",
7097   .num_srcs = 5,
7098   .src_components = {
7099      1, 1, 1, 1, 1
7100   },
7101   .has_dest = true,
7102   .dest_components = 1,
7103   .dest_bit_sizes = 0x0,
7104   .bit_size_src = -1,
7105   .num_indices = 1,
7106   .indices = {
7107      NIR_INTRINSIC_ACCESS,
7108   },
7109   .index_map = {
7110      [NIR_INTRINSIC_ACCESS] = 1,
7111    },
7112   .flags = 0,
7113},
7114{
7115   .name = "ssbo_atomic_exchange",
7116   .num_srcs = 3,
7117   .src_components = {
7118      -1, 1, 1
7119   },
7120   .has_dest = true,
7121   .dest_components = 1,
7122   .dest_bit_sizes = 0x0,
7123   .bit_size_src = -1,
7124   .num_indices = 1,
7125   .indices = {
7126      NIR_INTRINSIC_ACCESS,
7127   },
7128   .index_map = {
7129      [NIR_INTRINSIC_ACCESS] = 1,
7130    },
7131   .flags = 0,
7132},
7133{
7134   .name = "ssbo_atomic_exchange_ir3",
7135   .num_srcs = 4,
7136   .src_components = {
7137      1, 1, 1, 1
7138   },
7139   .has_dest = true,
7140   .dest_components = 1,
7141   .dest_bit_sizes = 0x0,
7142   .bit_size_src = -1,
7143   .num_indices = 1,
7144   .indices = {
7145      NIR_INTRINSIC_ACCESS,
7146   },
7147   .index_map = {
7148      [NIR_INTRINSIC_ACCESS] = 1,
7149    },
7150   .flags = 0,
7151},
7152{
7153   .name = "ssbo_atomic_fadd",
7154   .num_srcs = 3,
7155   .src_components = {
7156      -1, 1, 1
7157   },
7158   .has_dest = true,
7159   .dest_components = 1,
7160   .dest_bit_sizes = 0x0,
7161   .bit_size_src = -1,
7162   .num_indices = 1,
7163   .indices = {
7164      NIR_INTRINSIC_ACCESS,
7165   },
7166   .index_map = {
7167      [NIR_INTRINSIC_ACCESS] = 1,
7168    },
7169   .flags = 0,
7170},
7171{
7172   .name = "ssbo_atomic_fcomp_swap",
7173   .num_srcs = 4,
7174   .src_components = {
7175      -1, 1, 1, 1
7176   },
7177   .has_dest = true,
7178   .dest_components = 1,
7179   .dest_bit_sizes = 0x0,
7180   .bit_size_src = -1,
7181   .num_indices = 1,
7182   .indices = {
7183      NIR_INTRINSIC_ACCESS,
7184   },
7185   .index_map = {
7186      [NIR_INTRINSIC_ACCESS] = 1,
7187    },
7188   .flags = 0,
7189},
7190{
7191   .name = "ssbo_atomic_fmax",
7192   .num_srcs = 3,
7193   .src_components = {
7194      -1, 1, 1
7195   },
7196   .has_dest = true,
7197   .dest_components = 1,
7198   .dest_bit_sizes = 0x0,
7199   .bit_size_src = -1,
7200   .num_indices = 1,
7201   .indices = {
7202      NIR_INTRINSIC_ACCESS,
7203   },
7204   .index_map = {
7205      [NIR_INTRINSIC_ACCESS] = 1,
7206    },
7207   .flags = 0,
7208},
7209{
7210   .name = "ssbo_atomic_fmin",
7211   .num_srcs = 3,
7212   .src_components = {
7213      -1, 1, 1
7214   },
7215   .has_dest = true,
7216   .dest_components = 1,
7217   .dest_bit_sizes = 0x0,
7218   .bit_size_src = -1,
7219   .num_indices = 1,
7220   .indices = {
7221      NIR_INTRINSIC_ACCESS,
7222   },
7223   .index_map = {
7224      [NIR_INTRINSIC_ACCESS] = 1,
7225    },
7226   .flags = 0,
7227},
7228{
7229   .name = "ssbo_atomic_imax",
7230   .num_srcs = 3,
7231   .src_components = {
7232      -1, 1, 1
7233   },
7234   .has_dest = true,
7235   .dest_components = 1,
7236   .dest_bit_sizes = 0x0,
7237   .bit_size_src = -1,
7238   .num_indices = 1,
7239   .indices = {
7240      NIR_INTRINSIC_ACCESS,
7241   },
7242   .index_map = {
7243      [NIR_INTRINSIC_ACCESS] = 1,
7244    },
7245   .flags = 0,
7246},
7247{
7248   .name = "ssbo_atomic_imax_ir3",
7249   .num_srcs = 4,
7250   .src_components = {
7251      1, 1, 1, 1
7252   },
7253   .has_dest = true,
7254   .dest_components = 1,
7255   .dest_bit_sizes = 0x0,
7256   .bit_size_src = -1,
7257   .num_indices = 1,
7258   .indices = {
7259      NIR_INTRINSIC_ACCESS,
7260   },
7261   .index_map = {
7262      [NIR_INTRINSIC_ACCESS] = 1,
7263    },
7264   .flags = 0,
7265},
7266{
7267   .name = "ssbo_atomic_imin",
7268   .num_srcs = 3,
7269   .src_components = {
7270      -1, 1, 1
7271   },
7272   .has_dest = true,
7273   .dest_components = 1,
7274   .dest_bit_sizes = 0x0,
7275   .bit_size_src = -1,
7276   .num_indices = 1,
7277   .indices = {
7278      NIR_INTRINSIC_ACCESS,
7279   },
7280   .index_map = {
7281      [NIR_INTRINSIC_ACCESS] = 1,
7282    },
7283   .flags = 0,
7284},
7285{
7286   .name = "ssbo_atomic_imin_ir3",
7287   .num_srcs = 4,
7288   .src_components = {
7289      1, 1, 1, 1
7290   },
7291   .has_dest = true,
7292   .dest_components = 1,
7293   .dest_bit_sizes = 0x0,
7294   .bit_size_src = -1,
7295   .num_indices = 1,
7296   .indices = {
7297      NIR_INTRINSIC_ACCESS,
7298   },
7299   .index_map = {
7300      [NIR_INTRINSIC_ACCESS] = 1,
7301    },
7302   .flags = 0,
7303},
7304{
7305   .name = "ssbo_atomic_or",
7306   .num_srcs = 3,
7307   .src_components = {
7308      -1, 1, 1
7309   },
7310   .has_dest = true,
7311   .dest_components = 1,
7312   .dest_bit_sizes = 0x0,
7313   .bit_size_src = -1,
7314   .num_indices = 1,
7315   .indices = {
7316      NIR_INTRINSIC_ACCESS,
7317   },
7318   .index_map = {
7319      [NIR_INTRINSIC_ACCESS] = 1,
7320    },
7321   .flags = 0,
7322},
7323{
7324   .name = "ssbo_atomic_or_ir3",
7325   .num_srcs = 4,
7326   .src_components = {
7327      1, 1, 1, 1
7328   },
7329   .has_dest = true,
7330   .dest_components = 1,
7331   .dest_bit_sizes = 0x0,
7332   .bit_size_src = -1,
7333   .num_indices = 1,
7334   .indices = {
7335      NIR_INTRINSIC_ACCESS,
7336   },
7337   .index_map = {
7338      [NIR_INTRINSIC_ACCESS] = 1,
7339    },
7340   .flags = 0,
7341},
7342{
7343   .name = "ssbo_atomic_umax",
7344   .num_srcs = 3,
7345   .src_components = {
7346      -1, 1, 1
7347   },
7348   .has_dest = true,
7349   .dest_components = 1,
7350   .dest_bit_sizes = 0x0,
7351   .bit_size_src = -1,
7352   .num_indices = 1,
7353   .indices = {
7354      NIR_INTRINSIC_ACCESS,
7355   },
7356   .index_map = {
7357      [NIR_INTRINSIC_ACCESS] = 1,
7358    },
7359   .flags = 0,
7360},
7361{
7362   .name = "ssbo_atomic_umax_ir3",
7363   .num_srcs = 4,
7364   .src_components = {
7365      1, 1, 1, 1
7366   },
7367   .has_dest = true,
7368   .dest_components = 1,
7369   .dest_bit_sizes = 0x0,
7370   .bit_size_src = -1,
7371   .num_indices = 1,
7372   .indices = {
7373      NIR_INTRINSIC_ACCESS,
7374   },
7375   .index_map = {
7376      [NIR_INTRINSIC_ACCESS] = 1,
7377    },
7378   .flags = 0,
7379},
7380{
7381   .name = "ssbo_atomic_umin",
7382   .num_srcs = 3,
7383   .src_components = {
7384      -1, 1, 1
7385   },
7386   .has_dest = true,
7387   .dest_components = 1,
7388   .dest_bit_sizes = 0x0,
7389   .bit_size_src = -1,
7390   .num_indices = 1,
7391   .indices = {
7392      NIR_INTRINSIC_ACCESS,
7393   },
7394   .index_map = {
7395      [NIR_INTRINSIC_ACCESS] = 1,
7396    },
7397   .flags = 0,
7398},
7399{
7400   .name = "ssbo_atomic_umin_ir3",
7401   .num_srcs = 4,
7402   .src_components = {
7403      1, 1, 1, 1
7404   },
7405   .has_dest = true,
7406   .dest_components = 1,
7407   .dest_bit_sizes = 0x0,
7408   .bit_size_src = -1,
7409   .num_indices = 1,
7410   .indices = {
7411      NIR_INTRINSIC_ACCESS,
7412   },
7413   .index_map = {
7414      [NIR_INTRINSIC_ACCESS] = 1,
7415    },
7416   .flags = 0,
7417},
7418{
7419   .name = "ssbo_atomic_xor",
7420   .num_srcs = 3,
7421   .src_components = {
7422      -1, 1, 1
7423   },
7424   .has_dest = true,
7425   .dest_components = 1,
7426   .dest_bit_sizes = 0x0,
7427   .bit_size_src = -1,
7428   .num_indices = 1,
7429   .indices = {
7430      NIR_INTRINSIC_ACCESS,
7431   },
7432   .index_map = {
7433      [NIR_INTRINSIC_ACCESS] = 1,
7434    },
7435   .flags = 0,
7436},
7437{
7438   .name = "ssbo_atomic_xor_ir3",
7439   .num_srcs = 4,
7440   .src_components = {
7441      1, 1, 1, 1
7442   },
7443   .has_dest = true,
7444   .dest_components = 1,
7445   .dest_bit_sizes = 0x0,
7446   .bit_size_src = -1,
7447   .num_indices = 1,
7448   .indices = {
7449      NIR_INTRINSIC_ACCESS,
7450   },
7451   .index_map = {
7452      [NIR_INTRINSIC_ACCESS] = 1,
7453    },
7454   .flags = 0,
7455},
7456{
7457   .name = "store_buffer_amd",
7458   .num_srcs = 4,
7459   .src_components = {
7460      0, 4, 1, 1
7461   },
7462   .has_dest = false,
7463   .dest_components = 0,
7464   .dest_bit_sizes = 0x0,
7465   .bit_size_src = -1,
7466   .num_indices = 5,
7467   .indices = {
7468      NIR_INTRINSIC_BASE,
7469      NIR_INTRINSIC_WRITE_MASK,
7470      NIR_INTRINSIC_IS_SWIZZLED,
7471      NIR_INTRINSIC_SLC_AMD,
7472      NIR_INTRINSIC_MEMORY_MODES,
7473   },
7474   .index_map = {
7475      [NIR_INTRINSIC_BASE] = 1,
7476      [NIR_INTRINSIC_WRITE_MASK] = 2,
7477      [NIR_INTRINSIC_IS_SWIZZLED] = 3,
7478      [NIR_INTRINSIC_SLC_AMD] = 4,
7479      [NIR_INTRINSIC_MEMORY_MODES] = 5,
7480    },
7481   .flags = 0,
7482},
7483{
7484   .name = "store_combined_output_pan",
7485   .num_srcs = 4,
7486   .src_components = {
7487      0, 1, 1, 1
7488   },
7489   .has_dest = false,
7490   .dest_components = 0,
7491   .dest_bit_sizes = 0x0,
7492   .bit_size_src = -1,
7493   .num_indices = 3,
7494   .indices = {
7495      NIR_INTRINSIC_BASE,
7496      NIR_INTRINSIC_COMPONENT,
7497      NIR_INTRINSIC_SRC_TYPE,
7498   },
7499   .index_map = {
7500      [NIR_INTRINSIC_BASE] = 1,
7501      [NIR_INTRINSIC_COMPONENT] = 2,
7502      [NIR_INTRINSIC_SRC_TYPE] = 3,
7503    },
7504   .flags = 0,
7505},
7506{
7507   .name = "store_deref",
7508   .num_srcs = 2,
7509   .src_components = {
7510      -1, 0
7511   },
7512   .has_dest = false,
7513   .dest_components = 0,
7514   .dest_bit_sizes = 0x0,
7515   .bit_size_src = -1,
7516   .num_indices = 2,
7517   .indices = {
7518      NIR_INTRINSIC_WRITE_MASK,
7519      NIR_INTRINSIC_ACCESS,
7520   },
7521   .index_map = {
7522      [NIR_INTRINSIC_WRITE_MASK] = 1,
7523      [NIR_INTRINSIC_ACCESS] = 2,
7524    },
7525   .flags = 0,
7526},
7527{
7528   .name = "store_deref_block_intel",
7529   .num_srcs = 2,
7530   .src_components = {
7531      -1, 0
7532   },
7533   .has_dest = false,
7534   .dest_components = 0,
7535   .dest_bit_sizes = 0x0,
7536   .bit_size_src = -1,
7537   .num_indices = 2,
7538   .indices = {
7539      NIR_INTRINSIC_WRITE_MASK,
7540      NIR_INTRINSIC_ACCESS,
7541   },
7542   .index_map = {
7543      [NIR_INTRINSIC_WRITE_MASK] = 1,
7544      [NIR_INTRINSIC_ACCESS] = 2,
7545    },
7546   .flags = 0,
7547},
7548{
7549   .name = "store_global",
7550   .num_srcs = 2,
7551   .src_components = {
7552      0, 1
7553   },
7554   .has_dest = false,
7555   .dest_components = 0,
7556   .dest_bit_sizes = 0x0,
7557   .bit_size_src = -1,
7558   .num_indices = 4,
7559   .indices = {
7560      NIR_INTRINSIC_WRITE_MASK,
7561      NIR_INTRINSIC_ACCESS,
7562      NIR_INTRINSIC_ALIGN_MUL,
7563      NIR_INTRINSIC_ALIGN_OFFSET,
7564   },
7565   .index_map = {
7566      [NIR_INTRINSIC_WRITE_MASK] = 1,
7567      [NIR_INTRINSIC_ACCESS] = 2,
7568      [NIR_INTRINSIC_ALIGN_MUL] = 3,
7569      [NIR_INTRINSIC_ALIGN_OFFSET] = 4,
7570    },
7571   .flags = 0,
7572},
7573{
7574   .name = "store_global_block_intel",
7575   .num_srcs = 2,
7576   .src_components = {
7577      0, 1
7578   },
7579   .has_dest = false,
7580   .dest_components = 0,
7581   .dest_bit_sizes = 0x0,
7582   .bit_size_src = -1,
7583   .num_indices = 4,
7584   .indices = {
7585      NIR_INTRINSIC_WRITE_MASK,
7586      NIR_INTRINSIC_ACCESS,
7587      NIR_INTRINSIC_ALIGN_MUL,
7588      NIR_INTRINSIC_ALIGN_OFFSET,
7589   },
7590   .index_map = {
7591      [NIR_INTRINSIC_WRITE_MASK] = 1,
7592      [NIR_INTRINSIC_ACCESS] = 2,
7593      [NIR_INTRINSIC_ALIGN_MUL] = 3,
7594      [NIR_INTRINSIC_ALIGN_OFFSET] = 4,
7595    },
7596   .flags = 0,
7597},
7598{
7599   .name = "store_global_ir3",
7600   .num_srcs = 3,
7601   .src_components = {
7602      0, 2, 1
7603   },
7604   .has_dest = false,
7605   .dest_components = 0,
7606   .dest_bit_sizes = 0x0,
7607   .bit_size_src = -1,
7608   .num_indices = 3,
7609   .indices = {
7610      NIR_INTRINSIC_ACCESS,
7611      NIR_INTRINSIC_ALIGN_MUL,
7612      NIR_INTRINSIC_ALIGN_OFFSET,
7613   },
7614   .index_map = {
7615      [NIR_INTRINSIC_ACCESS] = 1,
7616      [NIR_INTRINSIC_ALIGN_MUL] = 2,
7617      [NIR_INTRINSIC_ALIGN_OFFSET] = 3,
7618    },
7619   .flags = 0,
7620},
7621{
7622   .name = "store_local_shared_r600",
7623   .num_srcs = 2,
7624   .src_components = {
7625      0, 1
7626   },
7627   .has_dest = false,
7628   .dest_components = 0,
7629   .dest_bit_sizes = 0x0,
7630   .bit_size_src = -1,
7631   .num_indices = 1,
7632   .indices = {
7633      NIR_INTRINSIC_WRITE_MASK,
7634   },
7635   .index_map = {
7636      [NIR_INTRINSIC_WRITE_MASK] = 1,
7637    },
7638   .flags = 0,
7639},
7640{
7641   .name = "store_output",
7642   .num_srcs = 2,
7643   .src_components = {
7644      0, 1
7645   },
7646   .has_dest = false,
7647   .dest_components = 0,
7648   .dest_bit_sizes = 0x0,
7649   .bit_size_src = -1,
7650   .num_indices = 5,
7651   .indices = {
7652      NIR_INTRINSIC_BASE,
7653      NIR_INTRINSIC_WRITE_MASK,
7654      NIR_INTRINSIC_COMPONENT,
7655      NIR_INTRINSIC_SRC_TYPE,
7656      NIR_INTRINSIC_IO_SEMANTICS,
7657   },
7658   .index_map = {
7659      [NIR_INTRINSIC_BASE] = 1,
7660      [NIR_INTRINSIC_WRITE_MASK] = 2,
7661      [NIR_INTRINSIC_COMPONENT] = 3,
7662      [NIR_INTRINSIC_SRC_TYPE] = 4,
7663      [NIR_INTRINSIC_IO_SEMANTICS] = 5,
7664    },
7665   .flags = 0,
7666},
7667{
7668   .name = "store_per_primitive_output",
7669   .num_srcs = 3,
7670   .src_components = {
7671      0, 1, 1
7672   },
7673   .has_dest = false,
7674   .dest_components = 0,
7675   .dest_bit_sizes = 0x0,
7676   .bit_size_src = -1,
7677   .num_indices = 5,
7678   .indices = {
7679      NIR_INTRINSIC_BASE,
7680      NIR_INTRINSIC_WRITE_MASK,
7681      NIR_INTRINSIC_COMPONENT,
7682      NIR_INTRINSIC_SRC_TYPE,
7683      NIR_INTRINSIC_IO_SEMANTICS,
7684   },
7685   .index_map = {
7686      [NIR_INTRINSIC_BASE] = 1,
7687      [NIR_INTRINSIC_WRITE_MASK] = 2,
7688      [NIR_INTRINSIC_COMPONENT] = 3,
7689      [NIR_INTRINSIC_SRC_TYPE] = 4,
7690      [NIR_INTRINSIC_IO_SEMANTICS] = 5,
7691    },
7692   .flags = 0,
7693},
7694{
7695   .name = "store_per_vertex_output",
7696   .num_srcs = 3,
7697   .src_components = {
7698      0, 1, 1
7699   },
7700   .has_dest = false,
7701   .dest_components = 0,
7702   .dest_bit_sizes = 0x0,
7703   .bit_size_src = -1,
7704   .num_indices = 5,
7705   .indices = {
7706      NIR_INTRINSIC_BASE,
7707      NIR_INTRINSIC_WRITE_MASK,
7708      NIR_INTRINSIC_COMPONENT,
7709      NIR_INTRINSIC_SRC_TYPE,
7710      NIR_INTRINSIC_IO_SEMANTICS,
7711   },
7712   .index_map = {
7713      [NIR_INTRINSIC_BASE] = 1,
7714      [NIR_INTRINSIC_WRITE_MASK] = 2,
7715      [NIR_INTRINSIC_COMPONENT] = 3,
7716      [NIR_INTRINSIC_SRC_TYPE] = 4,
7717      [NIR_INTRINSIC_IO_SEMANTICS] = 5,
7718    },
7719   .flags = 0,
7720},
7721{
7722   .name = "store_raw_output_pan",
7723   .num_srcs = 1,
7724   .src_components = {
7725      0
7726   },
7727   .has_dest = false,
7728   .dest_components = 0,
7729   .dest_bit_sizes = 0x0,
7730   .bit_size_src = -1,
7731   .num_indices = 0,
7732   .flags = 0,
7733},
7734{
7735   .name = "store_scratch",
7736   .num_srcs = 2,
7737   .src_components = {
7738      0, 1
7739   },
7740   .has_dest = false,
7741   .dest_components = 0,
7742   .dest_bit_sizes = 0x0,
7743   .bit_size_src = -1,
7744   .num_indices = 3,
7745   .indices = {
7746      NIR_INTRINSIC_ALIGN_MUL,
7747      NIR_INTRINSIC_ALIGN_OFFSET,
7748      NIR_INTRINSIC_WRITE_MASK,
7749   },
7750   .index_map = {
7751      [NIR_INTRINSIC_ALIGN_MUL] = 1,
7752      [NIR_INTRINSIC_ALIGN_OFFSET] = 2,
7753      [NIR_INTRINSIC_WRITE_MASK] = 3,
7754    },
7755   .flags = 0,
7756},
7757{
7758   .name = "store_scratch_dxil",
7759   .num_srcs = 2,
7760   .src_components = {
7761      1, 1
7762   },
7763   .has_dest = false,
7764   .dest_components = 0,
7765   .dest_bit_sizes = 0x0,
7766   .bit_size_src = -1,
7767   .num_indices = 0,
7768   .flags = 0,
7769},
7770{
7771   .name = "store_shared",
7772   .num_srcs = 2,
7773   .src_components = {
7774      0, 1
7775   },
7776   .has_dest = false,
7777   .dest_components = 0,
7778   .dest_bit_sizes = 0x0,
7779   .bit_size_src = -1,
7780   .num_indices = 4,
7781   .indices = {
7782      NIR_INTRINSIC_BASE,
7783      NIR_INTRINSIC_WRITE_MASK,
7784      NIR_INTRINSIC_ALIGN_MUL,
7785      NIR_INTRINSIC_ALIGN_OFFSET,
7786   },
7787   .index_map = {
7788      [NIR_INTRINSIC_BASE] = 1,
7789      [NIR_INTRINSIC_WRITE_MASK] = 2,
7790      [NIR_INTRINSIC_ALIGN_MUL] = 3,
7791      [NIR_INTRINSIC_ALIGN_OFFSET] = 4,
7792    },
7793   .flags = 0,
7794},
7795{
7796   .name = "store_shared_block_intel",
7797   .num_srcs = 2,
7798   .src_components = {
7799      0, 1
7800   },
7801   .has_dest = false,
7802   .dest_components = 0,
7803   .dest_bit_sizes = 0x0,
7804   .bit_size_src = -1,
7805   .num_indices = 4,
7806   .indices = {
7807      NIR_INTRINSIC_BASE,
7808      NIR_INTRINSIC_WRITE_MASK,
7809      NIR_INTRINSIC_ALIGN_MUL,
7810      NIR_INTRINSIC_ALIGN_OFFSET,
7811   },
7812   .index_map = {
7813      [NIR_INTRINSIC_BASE] = 1,
7814      [NIR_INTRINSIC_WRITE_MASK] = 2,
7815      [NIR_INTRINSIC_ALIGN_MUL] = 3,
7816      [NIR_INTRINSIC_ALIGN_OFFSET] = 4,
7817    },
7818   .flags = 0,
7819},
7820{
7821   .name = "store_shared_dxil",
7822   .num_srcs = 2,
7823   .src_components = {
7824      1, 1
7825   },
7826   .has_dest = false,
7827   .dest_components = 0,
7828   .dest_bit_sizes = 0x0,
7829   .bit_size_src = -1,
7830   .num_indices = 0,
7831   .flags = 0,
7832},
7833{
7834   .name = "store_shared_ir3",
7835   .num_srcs = 2,
7836   .src_components = {
7837      0, 1
7838   },
7839   .has_dest = false,
7840   .dest_components = 0,
7841   .dest_bit_sizes = 0x0,
7842   .bit_size_src = -1,
7843   .num_indices = 3,
7844   .indices = {
7845      NIR_INTRINSIC_BASE,
7846      NIR_INTRINSIC_ALIGN_MUL,
7847      NIR_INTRINSIC_ALIGN_OFFSET,
7848   },
7849   .index_map = {
7850      [NIR_INTRINSIC_BASE] = 1,
7851      [NIR_INTRINSIC_ALIGN_MUL] = 2,
7852      [NIR_INTRINSIC_ALIGN_OFFSET] = 3,
7853    },
7854   .flags = 0,
7855},
7856{
7857   .name = "store_shared_masked_dxil",
7858   .num_srcs = 3,
7859   .src_components = {
7860      1, 1, 1
7861   },
7862   .has_dest = false,
7863   .dest_components = 0,
7864   .dest_bit_sizes = 0x0,
7865   .bit_size_src = -1,
7866   .num_indices = 0,
7867   .flags = 0,
7868},
7869{
7870   .name = "store_ssbo",
7871   .num_srcs = 3,
7872   .src_components = {
7873      0, -1, 1
7874   },
7875   .has_dest = false,
7876   .dest_components = 0,
7877   .dest_bit_sizes = 0x0,
7878   .bit_size_src = -1,
7879   .num_indices = 4,
7880   .indices = {
7881      NIR_INTRINSIC_WRITE_MASK,
7882      NIR_INTRINSIC_ACCESS,
7883      NIR_INTRINSIC_ALIGN_MUL,
7884      NIR_INTRINSIC_ALIGN_OFFSET,
7885   },
7886   .index_map = {
7887      [NIR_INTRINSIC_WRITE_MASK] = 1,
7888      [NIR_INTRINSIC_ACCESS] = 2,
7889      [NIR_INTRINSIC_ALIGN_MUL] = 3,
7890      [NIR_INTRINSIC_ALIGN_OFFSET] = 4,
7891    },
7892   .flags = 0,
7893},
7894{
7895   .name = "store_ssbo_block_intel",
7896   .num_srcs = 3,
7897   .src_components = {
7898      0, -1, 1
7899   },
7900   .has_dest = false,
7901   .dest_components = 0,
7902   .dest_bit_sizes = 0x0,
7903   .bit_size_src = -1,
7904   .num_indices = 4,
7905   .indices = {
7906      NIR_INTRINSIC_WRITE_MASK,
7907      NIR_INTRINSIC_ACCESS,
7908      NIR_INTRINSIC_ALIGN_MUL,
7909      NIR_INTRINSIC_ALIGN_OFFSET,
7910   },
7911   .index_map = {
7912      [NIR_INTRINSIC_WRITE_MASK] = 1,
7913      [NIR_INTRINSIC_ACCESS] = 2,
7914      [NIR_INTRINSIC_ALIGN_MUL] = 3,
7915      [NIR_INTRINSIC_ALIGN_OFFSET] = 4,
7916    },
7917   .flags = 0,
7918},
7919{
7920   .name = "store_ssbo_ir3",
7921   .num_srcs = 4,
7922   .src_components = {
7923      0, 1, 1, 1
7924   },
7925   .has_dest = false,
7926   .dest_components = 0,
7927   .dest_bit_sizes = 0x0,
7928   .bit_size_src = -1,
7929   .num_indices = 4,
7930   .indices = {
7931      NIR_INTRINSIC_WRITE_MASK,
7932      NIR_INTRINSIC_ACCESS,
7933      NIR_INTRINSIC_ALIGN_MUL,
7934      NIR_INTRINSIC_ALIGN_OFFSET,
7935   },
7936   .index_map = {
7937      [NIR_INTRINSIC_WRITE_MASK] = 1,
7938      [NIR_INTRINSIC_ACCESS] = 2,
7939      [NIR_INTRINSIC_ALIGN_MUL] = 3,
7940      [NIR_INTRINSIC_ALIGN_OFFSET] = 4,
7941    },
7942   .flags = 0,
7943},
7944{
7945   .name = "store_ssbo_masked_dxil",
7946   .num_srcs = 4,
7947   .src_components = {
7948      1, 1, 1, 1
7949   },
7950   .has_dest = false,
7951   .dest_components = 0,
7952   .dest_bit_sizes = 0x0,
7953   .bit_size_src = -1,
7954   .num_indices = 0,
7955   .flags = 0,
7956},
7957{
7958   .name = "store_tf_r600",
7959   .num_srcs = 1,
7960   .src_components = {
7961      0
7962   },
7963   .has_dest = false,
7964   .dest_components = 0,
7965   .dest_bit_sizes = 0x0,
7966   .bit_size_src = -1,
7967   .num_indices = 0,
7968   .flags = 0,
7969},
7970{
7971   .name = "store_tlb_sample_color_v3d",
7972   .num_srcs = 2,
7973   .src_components = {
7974      0, 1
7975   },
7976   .has_dest = false,
7977   .dest_components = 0,
7978   .dest_bit_sizes = 0x0,
7979   .bit_size_src = -1,
7980   .num_indices = 3,
7981   .indices = {
7982      NIR_INTRINSIC_BASE,
7983      NIR_INTRINSIC_COMPONENT,
7984      NIR_INTRINSIC_SRC_TYPE,
7985   },
7986   .index_map = {
7987      [NIR_INTRINSIC_BASE] = 1,
7988      [NIR_INTRINSIC_COMPONENT] = 2,
7989      [NIR_INTRINSIC_SRC_TYPE] = 3,
7990    },
7991   .flags = 0,
7992},
7993{
7994   .name = "terminate",
7995   .num_srcs = 0,
7996   .has_dest = false,
7997   .dest_components = 0,
7998   .dest_bit_sizes = 0x0,
7999   .bit_size_src = -1,
8000   .num_indices = 0,
8001   .flags = 0,
8002},
8003{
8004   .name = "terminate_if",
8005   .num_srcs = 1,
8006   .src_components = {
8007      1
8008   },
8009   .has_dest = false,
8010   .dest_components = 0,
8011   .dest_bit_sizes = 0x0,
8012   .bit_size_src = -1,
8013   .num_indices = 0,
8014   .flags = 0,
8015},
8016{
8017   .name = "terminate_ray",
8018   .num_srcs = 0,
8019   .has_dest = false,
8020   .dest_components = 0,
8021   .dest_bit_sizes = 0x0,
8022   .bit_size_src = -1,
8023   .num_indices = 0,
8024   .flags = 0,
8025},
8026{
8027   .name = "trace_ray",
8028   .num_srcs = 11,
8029   .src_components = {
8030      -1, 1, 1, 1, 1, 1, 3, 1, 3, 1, -1
8031   },
8032   .has_dest = false,
8033   .dest_components = 0,
8034   .dest_bit_sizes = 0x0,
8035   .bit_size_src = -1,
8036   .num_indices = 0,
8037   .flags = 0,
8038},
8039{
8040   .name = "trace_ray_commit_intel",
8041   .num_srcs = 0,
8042   .has_dest = false,
8043   .dest_components = 0,
8044   .dest_bit_sizes = 0x0,
8045   .bit_size_src = -1,
8046   .num_indices = 0,
8047   .flags = 0,
8048},
8049{
8050   .name = "trace_ray_continue_intel",
8051   .num_srcs = 0,
8052   .has_dest = false,
8053   .dest_components = 0,
8054   .dest_bit_sizes = 0x0,
8055   .bit_size_src = -1,
8056   .num_indices = 0,
8057   .flags = 0,
8058},
8059{
8060   .name = "trace_ray_initial_intel",
8061   .num_srcs = 0,
8062   .has_dest = false,
8063   .dest_components = 0,
8064   .dest_bit_sizes = 0x0,
8065   .bit_size_src = -1,
8066   .num_indices = 0,
8067   .flags = 0,
8068},
8069{
8070   .name = "vote_all",
8071   .num_srcs = 1,
8072   .src_components = {
8073      1
8074   },
8075   .has_dest = true,
8076   .dest_components = 1,
8077   .dest_bit_sizes = 0x0,
8078   .bit_size_src = -1,
8079   .num_indices = 0,
8080   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
8081},
8082{
8083   .name = "vote_any",
8084   .num_srcs = 1,
8085   .src_components = {
8086      1
8087   },
8088   .has_dest = true,
8089   .dest_components = 1,
8090   .dest_bit_sizes = 0x0,
8091   .bit_size_src = -1,
8092   .num_indices = 0,
8093   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
8094},
8095{
8096   .name = "vote_feq",
8097   .num_srcs = 1,
8098   .src_components = {
8099      0
8100   },
8101   .has_dest = true,
8102   .dest_components = 1,
8103   .dest_bit_sizes = 0x0,
8104   .bit_size_src = -1,
8105   .num_indices = 0,
8106   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
8107},
8108{
8109   .name = "vote_ieq",
8110   .num_srcs = 1,
8111   .src_components = {
8112      0
8113   },
8114   .has_dest = true,
8115   .dest_components = 1,
8116   .dest_bit_sizes = 0x0,
8117   .bit_size_src = -1,
8118   .num_indices = 0,
8119   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
8120},
8121{
8122   .name = "vulkan_resource_index",
8123   .num_srcs = 1,
8124   .src_components = {
8125      1
8126   },
8127   .has_dest = true,
8128   .dest_components = 0,
8129   .dest_bit_sizes = 0x0,
8130   .bit_size_src = -1,
8131   .num_indices = 3,
8132   .indices = {
8133      NIR_INTRINSIC_DESC_SET,
8134      NIR_INTRINSIC_BINDING,
8135      NIR_INTRINSIC_DESC_TYPE,
8136   },
8137   .index_map = {
8138      [NIR_INTRINSIC_DESC_SET] = 1,
8139      [NIR_INTRINSIC_BINDING] = 2,
8140      [NIR_INTRINSIC_DESC_TYPE] = 3,
8141    },
8142   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
8143},
8144{
8145   .name = "vulkan_resource_reindex",
8146   .num_srcs = 2,
8147   .src_components = {
8148      0, 1
8149   },
8150   .has_dest = true,
8151   .dest_components = 0,
8152   .dest_bit_sizes = 0x0,
8153   .bit_size_src = -1,
8154   .num_indices = 1,
8155   .indices = {
8156      NIR_INTRINSIC_DESC_TYPE,
8157   },
8158   .index_map = {
8159      [NIR_INTRINSIC_DESC_TYPE] = 1,
8160    },
8161   .flags = NIR_INTRINSIC_CAN_ELIMINATE | NIR_INTRINSIC_CAN_REORDER,
8162},
8163{
8164   .name = "write_invocation_amd",
8165   .num_srcs = 3,
8166   .src_components = {
8167      0, 0, 1
8168   },
8169   .has_dest = true,
8170   .dest_components = 0,
8171   .dest_bit_sizes = 0x0,
8172   .bit_size_src = 0,
8173   .num_indices = 0,
8174   .flags = NIR_INTRINSIC_CAN_ELIMINATE,
8175},
8176};
8177
8178const char *nir_intrinsic_index_names[NIR_INTRINSIC_NUM_INDEX_FLAGS] = {
8179   "base",
8180   "write_mask",
8181   "stream_id",
8182   "ucp_id",
8183   "range_base",
8184   "range",
8185   "desc_set",
8186   "binding",
8187   "component",
8188   "column",
8189   "interp_mode",
8190   "reduction_op",
8191   "cluster_size",
8192   "param_idx",
8193   "image_dim",
8194   "image_array",
8195   "format",
8196   "access",
8197   "call_idx",
8198   "stack_size",
8199   "align_mul",
8200   "align_offset",
8201   "desc_type",
8202   "src_type",
8203   "dest_type",
8204   "swizzle_mask",
8205   "is_swizzled",
8206   "slc_amd",
8207   "dst_access",
8208   "src_access",
8209   "driver_location",
8210   "memory_semantics",
8211   "memory_modes",
8212   "memory_scope",
8213   "execution_scope",
8214   "io_semantics",
8215   "rounding_mode",
8216   "saturate",
8217};
8218