2 * Copyright © 2015 Intel Corporation
4 * Permission is hereby granted, free of charge, to any person obtaining a
5 * copy of this software and associated documentation files (the "Software"),
6 * to deal in the Software without restriction, including without limitation
7 * the rights to use, copy, modify, merge, publish, distribute, sublicense,
8 * and/or sell copies of the Software, and to permit persons to whom the
9 * Software is furnished to do so, subject to the following conditions:
11 * The above copyright notice and this permission notice (including the next
12 * paragraph) shall be included in all copies or substantial portions of the
15 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
16 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
17 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
18 * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
19 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
20 * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
24 * Jason Ekstrand (jason@jlekstrand.net)
28 #include "vtn_private.h"
29 #include "nir/nir_vla.h"
30 #include "nir/nir_control_flow.h"
31 #include "nir/nir_constant_expressions.h"
32 #include "spirv_info.h"
34 struct spec_constant_value
{
43 _vtn_warn(const char *file
, int line
, const char *msg
, ...)
49 formatted
= ralloc_vasprintf(NULL
, msg
, args
);
52 fprintf(stderr
, "%s:%d WARNING: %s\n", file
, line
, formatted
);
54 ralloc_free(formatted
);
57 static struct vtn_ssa_value
*
58 vtn_undef_ssa_value(struct vtn_builder
*b
, const struct glsl_type
*type
)
60 struct vtn_ssa_value
*val
= rzalloc(b
, struct vtn_ssa_value
);
63 if (glsl_type_is_vector_or_scalar(type
)) {
64 unsigned num_components
= glsl_get_vector_elements(val
->type
);
65 unsigned bit_size
= glsl_get_bit_size(val
->type
);
66 val
->def
= nir_ssa_undef(&b
->nb
, num_components
, bit_size
);
68 unsigned elems
= glsl_get_length(val
->type
);
69 val
->elems
= ralloc_array(b
, struct vtn_ssa_value
*, elems
);
70 if (glsl_type_is_matrix(type
)) {
71 const struct glsl_type
*elem_type
=
72 glsl_vector_type(glsl_get_base_type(type
),
73 glsl_get_vector_elements(type
));
75 for (unsigned i
= 0; i
< elems
; i
++)
76 val
->elems
[i
] = vtn_undef_ssa_value(b
, elem_type
);
77 } else if (glsl_type_is_array(type
)) {
78 const struct glsl_type
*elem_type
= glsl_get_array_element(type
);
79 for (unsigned i
= 0; i
< elems
; i
++)
80 val
->elems
[i
] = vtn_undef_ssa_value(b
, elem_type
);
82 for (unsigned i
= 0; i
< elems
; i
++) {
83 const struct glsl_type
*elem_type
= glsl_get_struct_field(type
, i
);
84 val
->elems
[i
] = vtn_undef_ssa_value(b
, elem_type
);
92 static struct vtn_ssa_value
*
93 vtn_const_ssa_value(struct vtn_builder
*b
, nir_constant
*constant
,
94 const struct glsl_type
*type
)
96 struct hash_entry
*entry
= _mesa_hash_table_search(b
->const_table
, constant
);
101 struct vtn_ssa_value
*val
= rzalloc(b
, struct vtn_ssa_value
);
104 switch (glsl_get_base_type(type
)) {
107 case GLSL_TYPE_INT64
:
108 case GLSL_TYPE_UINT64
:
110 case GLSL_TYPE_FLOAT
:
111 case GLSL_TYPE_DOUBLE
: {
112 int bit_size
= glsl_get_bit_size(type
);
113 if (glsl_type_is_vector_or_scalar(type
)) {
114 unsigned num_components
= glsl_get_vector_elements(val
->type
);
115 nir_load_const_instr
*load
=
116 nir_load_const_instr_create(b
->shader
, num_components
, bit_size
);
118 load
->value
= constant
->values
[0];
120 nir_instr_insert_before_cf_list(&b
->impl
->body
, &load
->instr
);
121 val
->def
= &load
->def
;
123 assert(glsl_type_is_matrix(type
));
124 unsigned rows
= glsl_get_vector_elements(val
->type
);
125 unsigned columns
= glsl_get_matrix_columns(val
->type
);
126 val
->elems
= ralloc_array(b
, struct vtn_ssa_value
*, columns
);
128 for (unsigned i
= 0; i
< columns
; i
++) {
129 struct vtn_ssa_value
*col_val
= rzalloc(b
, struct vtn_ssa_value
);
130 col_val
->type
= glsl_get_column_type(val
->type
);
131 nir_load_const_instr
*load
=
132 nir_load_const_instr_create(b
->shader
, rows
, bit_size
);
134 load
->value
= constant
->values
[i
];
136 nir_instr_insert_before_cf_list(&b
->impl
->body
, &load
->instr
);
137 col_val
->def
= &load
->def
;
139 val
->elems
[i
] = col_val
;
145 case GLSL_TYPE_ARRAY
: {
146 unsigned elems
= glsl_get_length(val
->type
);
147 val
->elems
= ralloc_array(b
, struct vtn_ssa_value
*, elems
);
148 const struct glsl_type
*elem_type
= glsl_get_array_element(val
->type
);
149 for (unsigned i
= 0; i
< elems
; i
++)
150 val
->elems
[i
] = vtn_const_ssa_value(b
, constant
->elements
[i
],
155 case GLSL_TYPE_STRUCT
: {
156 unsigned elems
= glsl_get_length(val
->type
);
157 val
->elems
= ralloc_array(b
, struct vtn_ssa_value
*, elems
);
158 for (unsigned i
= 0; i
< elems
; i
++) {
159 const struct glsl_type
*elem_type
=
160 glsl_get_struct_field(val
->type
, i
);
161 val
->elems
[i
] = vtn_const_ssa_value(b
, constant
->elements
[i
],
168 unreachable("bad constant type");
174 struct vtn_ssa_value
*
175 vtn_ssa_value(struct vtn_builder
*b
, uint32_t value_id
)
177 struct vtn_value
*val
= vtn_untyped_value(b
, value_id
);
178 switch (val
->value_type
) {
179 case vtn_value_type_undef
:
180 return vtn_undef_ssa_value(b
, val
->type
->type
);
182 case vtn_value_type_constant
:
183 return vtn_const_ssa_value(b
, val
->constant
, val
->const_type
);
185 case vtn_value_type_ssa
:
188 case vtn_value_type_pointer
:
189 /* This is needed for function parameters */
190 return vtn_variable_load(b
, val
->pointer
);
193 unreachable("Invalid type for an SSA value");
198 vtn_string_literal(struct vtn_builder
*b
, const uint32_t *words
,
199 unsigned word_count
, unsigned *words_used
)
201 char *dup
= ralloc_strndup(b
, (char *)words
, word_count
* sizeof(*words
));
203 /* Ammount of space taken by the string (including the null) */
204 unsigned len
= strlen(dup
) + 1;
205 *words_used
= DIV_ROUND_UP(len
, sizeof(*words
));
211 vtn_foreach_instruction(struct vtn_builder
*b
, const uint32_t *start
,
212 const uint32_t *end
, vtn_instruction_handler handler
)
218 const uint32_t *w
= start
;
220 SpvOp opcode
= w
[0] & SpvOpCodeMask
;
221 unsigned count
= w
[0] >> SpvWordCountShift
;
222 assert(count
>= 1 && w
+ count
<= end
);
226 break; /* Do nothing */
229 b
->file
= vtn_value(b
, w
[1], vtn_value_type_string
)->str
;
241 if (!handler(b
, opcode
, w
, count
))
253 vtn_handle_extension(struct vtn_builder
*b
, SpvOp opcode
,
254 const uint32_t *w
, unsigned count
)
257 case SpvOpExtInstImport
: {
258 struct vtn_value
*val
= vtn_push_value(b
, w
[1], vtn_value_type_extension
);
259 if (strcmp((const char *)&w
[2], "GLSL.std.450") == 0) {
260 val
->ext_handler
= vtn_handle_glsl450_instruction
;
262 assert(!"Unsupported extension");
268 struct vtn_value
*val
= vtn_value(b
, w
[3], vtn_value_type_extension
);
269 bool handled
= val
->ext_handler(b
, w
[4], w
, count
);
276 unreachable("Unhandled opcode");
281 _foreach_decoration_helper(struct vtn_builder
*b
,
282 struct vtn_value
*base_value
,
284 struct vtn_value
*value
,
285 vtn_decoration_foreach_cb cb
, void *data
)
287 for (struct vtn_decoration
*dec
= value
->decoration
; dec
; dec
= dec
->next
) {
289 if (dec
->scope
== VTN_DEC_DECORATION
) {
290 member
= parent_member
;
291 } else if (dec
->scope
>= VTN_DEC_STRUCT_MEMBER0
) {
292 assert(parent_member
== -1);
293 member
= dec
->scope
- VTN_DEC_STRUCT_MEMBER0
;
295 /* Not a decoration */
300 assert(dec
->group
->value_type
== vtn_value_type_decoration_group
);
301 _foreach_decoration_helper(b
, base_value
, member
, dec
->group
,
304 cb(b
, base_value
, member
, dec
, data
);
309 /** Iterates (recursively if needed) over all of the decorations on a value
311 * This function iterates over all of the decorations applied to a given
312 * value. If it encounters a decoration group, it recurses into the group
313 * and iterates over all of those decorations as well.
316 vtn_foreach_decoration(struct vtn_builder
*b
, struct vtn_value
*value
,
317 vtn_decoration_foreach_cb cb
, void *data
)
319 _foreach_decoration_helper(b
, value
, -1, value
, cb
, data
);
323 vtn_foreach_execution_mode(struct vtn_builder
*b
, struct vtn_value
*value
,
324 vtn_execution_mode_foreach_cb cb
, void *data
)
326 for (struct vtn_decoration
*dec
= value
->decoration
; dec
; dec
= dec
->next
) {
327 if (dec
->scope
!= VTN_DEC_EXECUTION_MODE
)
330 assert(dec
->group
== NULL
);
331 cb(b
, value
, dec
, data
);
336 vtn_handle_decoration(struct vtn_builder
*b
, SpvOp opcode
,
337 const uint32_t *w
, unsigned count
)
339 const uint32_t *w_end
= w
+ count
;
340 const uint32_t target
= w
[1];
344 case SpvOpDecorationGroup
:
345 vtn_push_value(b
, target
, vtn_value_type_decoration_group
);
349 case SpvOpMemberDecorate
:
350 case SpvOpExecutionMode
: {
351 struct vtn_value
*val
= &b
->values
[target
];
353 struct vtn_decoration
*dec
= rzalloc(b
, struct vtn_decoration
);
356 dec
->scope
= VTN_DEC_DECORATION
;
358 case SpvOpMemberDecorate
:
359 dec
->scope
= VTN_DEC_STRUCT_MEMBER0
+ *(w
++);
361 case SpvOpExecutionMode
:
362 dec
->scope
= VTN_DEC_EXECUTION_MODE
;
365 unreachable("Invalid decoration opcode");
367 dec
->decoration
= *(w
++);
370 /* Link into the list */
371 dec
->next
= val
->decoration
;
372 val
->decoration
= dec
;
376 case SpvOpGroupMemberDecorate
:
377 case SpvOpGroupDecorate
: {
378 struct vtn_value
*group
=
379 vtn_value(b
, target
, vtn_value_type_decoration_group
);
381 for (; w
< w_end
; w
++) {
382 struct vtn_value
*val
= vtn_untyped_value(b
, *w
);
383 struct vtn_decoration
*dec
= rzalloc(b
, struct vtn_decoration
);
386 if (opcode
== SpvOpGroupDecorate
) {
387 dec
->scope
= VTN_DEC_DECORATION
;
389 dec
->scope
= VTN_DEC_STRUCT_MEMBER0
+ *(++w
);
392 /* Link into the list */
393 dec
->next
= val
->decoration
;
394 val
->decoration
= dec
;
400 unreachable("Unhandled opcode");
404 struct member_decoration_ctx
{
406 struct glsl_struct_field
*fields
;
407 struct vtn_type
*type
;
410 /* does a shallow copy of a vtn_type */
412 static struct vtn_type
*
413 vtn_type_copy(struct vtn_builder
*b
, struct vtn_type
*src
)
415 struct vtn_type
*dest
= ralloc(b
, struct vtn_type
);
418 if (!glsl_type_is_scalar(src
->type
)) {
419 switch (glsl_get_base_type(src
->type
)) {
422 case GLSL_TYPE_INT64
:
423 case GLSL_TYPE_UINT64
:
425 case GLSL_TYPE_FLOAT
:
426 case GLSL_TYPE_DOUBLE
:
427 case GLSL_TYPE_ARRAY
:
430 case GLSL_TYPE_STRUCT
: {
431 unsigned elems
= glsl_get_length(src
->type
);
433 dest
->members
= ralloc_array(b
, struct vtn_type
*, elems
);
434 memcpy(dest
->members
, src
->members
, elems
* sizeof(struct vtn_type
*));
436 dest
->offsets
= ralloc_array(b
, unsigned, elems
);
437 memcpy(dest
->offsets
, src
->offsets
, elems
* sizeof(unsigned));
442 unreachable("unhandled type");
449 static struct vtn_type
*
450 mutable_matrix_member(struct vtn_builder
*b
, struct vtn_type
*type
, int member
)
452 type
->members
[member
] = vtn_type_copy(b
, type
->members
[member
]);
453 type
= type
->members
[member
];
455 /* We may have an array of matrices.... Oh, joy! */
456 while (glsl_type_is_array(type
->type
)) {
457 type
->array_element
= vtn_type_copy(b
, type
->array_element
);
458 type
= type
->array_element
;
461 assert(glsl_type_is_matrix(type
->type
));
467 struct_member_decoration_cb(struct vtn_builder
*b
,
468 struct vtn_value
*val
, int member
,
469 const struct vtn_decoration
*dec
, void *void_ctx
)
471 struct member_decoration_ctx
*ctx
= void_ctx
;
476 assert(member
< ctx
->num_fields
);
478 switch (dec
->decoration
) {
479 case SpvDecorationNonWritable
:
480 case SpvDecorationNonReadable
:
481 case SpvDecorationRelaxedPrecision
:
482 case SpvDecorationVolatile
:
483 case SpvDecorationCoherent
:
484 case SpvDecorationUniform
:
485 break; /* FIXME: Do nothing with this for now. */
486 case SpvDecorationNoPerspective
:
487 ctx
->fields
[member
].interpolation
= INTERP_MODE_NOPERSPECTIVE
;
489 case SpvDecorationFlat
:
490 ctx
->fields
[member
].interpolation
= INTERP_MODE_FLAT
;
492 case SpvDecorationCentroid
:
493 ctx
->fields
[member
].centroid
= true;
495 case SpvDecorationSample
:
496 ctx
->fields
[member
].sample
= true;
498 case SpvDecorationStream
:
499 /* Vulkan only allows one GS stream */
500 assert(dec
->literals
[0] == 0);
502 case SpvDecorationLocation
:
503 ctx
->fields
[member
].location
= dec
->literals
[0];
505 case SpvDecorationComponent
:
506 break; /* FIXME: What should we do with these? */
507 case SpvDecorationBuiltIn
:
508 ctx
->type
->members
[member
] = vtn_type_copy(b
, ctx
->type
->members
[member
]);
509 ctx
->type
->members
[member
]->is_builtin
= true;
510 ctx
->type
->members
[member
]->builtin
= dec
->literals
[0];
511 ctx
->type
->builtin_block
= true;
513 case SpvDecorationOffset
:
514 ctx
->type
->offsets
[member
] = dec
->literals
[0];
516 case SpvDecorationMatrixStride
:
517 /* Handled as a second pass */
519 case SpvDecorationColMajor
:
520 break; /* Nothing to do here. Column-major is the default. */
521 case SpvDecorationRowMajor
:
522 mutable_matrix_member(b
, ctx
->type
, member
)->row_major
= true;
525 case SpvDecorationPatch
:
528 case SpvDecorationSpecId
:
529 case SpvDecorationBlock
:
530 case SpvDecorationBufferBlock
:
531 case SpvDecorationArrayStride
:
532 case SpvDecorationGLSLShared
:
533 case SpvDecorationGLSLPacked
:
534 case SpvDecorationInvariant
:
535 case SpvDecorationRestrict
:
536 case SpvDecorationAliased
:
537 case SpvDecorationConstant
:
538 case SpvDecorationIndex
:
539 case SpvDecorationBinding
:
540 case SpvDecorationDescriptorSet
:
541 case SpvDecorationLinkageAttributes
:
542 case SpvDecorationNoContraction
:
543 case SpvDecorationInputAttachmentIndex
:
544 vtn_warn("Decoration not allowed on struct members: %s",
545 spirv_decoration_to_string(dec
->decoration
));
548 case SpvDecorationXfbBuffer
:
549 case SpvDecorationXfbStride
:
550 vtn_warn("Vulkan does not have transform feedback");
553 case SpvDecorationCPacked
:
554 case SpvDecorationSaturatedConversion
:
555 case SpvDecorationFuncParamAttr
:
556 case SpvDecorationFPRoundingMode
:
557 case SpvDecorationFPFastMathMode
:
558 case SpvDecorationAlignment
:
559 vtn_warn("Decoration only allowed for CL-style kernels: %s",
560 spirv_decoration_to_string(dec
->decoration
));
564 unreachable("Unhandled decoration");
568 /* Matrix strides are handled as a separate pass because we need to know
569 * whether the matrix is row-major or not first.
572 struct_member_matrix_stride_cb(struct vtn_builder
*b
,
573 struct vtn_value
*val
, int member
,
574 const struct vtn_decoration
*dec
,
577 if (dec
->decoration
!= SpvDecorationMatrixStride
)
581 struct member_decoration_ctx
*ctx
= void_ctx
;
583 struct vtn_type
*mat_type
= mutable_matrix_member(b
, ctx
->type
, member
);
584 if (mat_type
->row_major
) {
585 mat_type
->array_element
= vtn_type_copy(b
, mat_type
->array_element
);
586 mat_type
->stride
= mat_type
->array_element
->stride
;
587 mat_type
->array_element
->stride
= dec
->literals
[0];
589 assert(mat_type
->array_element
->stride
> 0);
590 mat_type
->stride
= dec
->literals
[0];
595 type_decoration_cb(struct vtn_builder
*b
,
596 struct vtn_value
*val
, int member
,
597 const struct vtn_decoration
*dec
, void *ctx
)
599 struct vtn_type
*type
= val
->type
;
604 switch (dec
->decoration
) {
605 case SpvDecorationArrayStride
:
606 type
->stride
= dec
->literals
[0];
608 case SpvDecorationBlock
:
611 case SpvDecorationBufferBlock
:
612 type
->buffer_block
= true;
614 case SpvDecorationGLSLShared
:
615 case SpvDecorationGLSLPacked
:
616 /* Ignore these, since we get explicit offsets anyways */
619 case SpvDecorationRowMajor
:
620 case SpvDecorationColMajor
:
621 case SpvDecorationMatrixStride
:
622 case SpvDecorationBuiltIn
:
623 case SpvDecorationNoPerspective
:
624 case SpvDecorationFlat
:
625 case SpvDecorationPatch
:
626 case SpvDecorationCentroid
:
627 case SpvDecorationSample
:
628 case SpvDecorationVolatile
:
629 case SpvDecorationCoherent
:
630 case SpvDecorationNonWritable
:
631 case SpvDecorationNonReadable
:
632 case SpvDecorationUniform
:
633 case SpvDecorationStream
:
634 case SpvDecorationLocation
:
635 case SpvDecorationComponent
:
636 case SpvDecorationOffset
:
637 case SpvDecorationXfbBuffer
:
638 case SpvDecorationXfbStride
:
639 vtn_warn("Decoration only allowed for struct members: %s",
640 spirv_decoration_to_string(dec
->decoration
));
643 case SpvDecorationRelaxedPrecision
:
644 case SpvDecorationSpecId
:
645 case SpvDecorationInvariant
:
646 case SpvDecorationRestrict
:
647 case SpvDecorationAliased
:
648 case SpvDecorationConstant
:
649 case SpvDecorationIndex
:
650 case SpvDecorationBinding
:
651 case SpvDecorationDescriptorSet
:
652 case SpvDecorationLinkageAttributes
:
653 case SpvDecorationNoContraction
:
654 case SpvDecorationInputAttachmentIndex
:
655 vtn_warn("Decoration not allowed on types: %s",
656 spirv_decoration_to_string(dec
->decoration
));
659 case SpvDecorationCPacked
:
660 case SpvDecorationSaturatedConversion
:
661 case SpvDecorationFuncParamAttr
:
662 case SpvDecorationFPRoundingMode
:
663 case SpvDecorationFPFastMathMode
:
664 case SpvDecorationAlignment
:
665 vtn_warn("Decoration only allowed for CL-style kernels: %s",
666 spirv_decoration_to_string(dec
->decoration
));
670 unreachable("Unhandled decoration");
675 translate_image_format(SpvImageFormat format
)
678 case SpvImageFormatUnknown
: return 0; /* GL_NONE */
679 case SpvImageFormatRgba32f
: return 0x8814; /* GL_RGBA32F */
680 case SpvImageFormatRgba16f
: return 0x881A; /* GL_RGBA16F */
681 case SpvImageFormatR32f
: return 0x822E; /* GL_R32F */
682 case SpvImageFormatRgba8
: return 0x8058; /* GL_RGBA8 */
683 case SpvImageFormatRgba8Snorm
: return 0x8F97; /* GL_RGBA8_SNORM */
684 case SpvImageFormatRg32f
: return 0x8230; /* GL_RG32F */
685 case SpvImageFormatRg16f
: return 0x822F; /* GL_RG16F */
686 case SpvImageFormatR11fG11fB10f
: return 0x8C3A; /* GL_R11F_G11F_B10F */
687 case SpvImageFormatR16f
: return 0x822D; /* GL_R16F */
688 case SpvImageFormatRgba16
: return 0x805B; /* GL_RGBA16 */
689 case SpvImageFormatRgb10A2
: return 0x8059; /* GL_RGB10_A2 */
690 case SpvImageFormatRg16
: return 0x822C; /* GL_RG16 */
691 case SpvImageFormatRg8
: return 0x822B; /* GL_RG8 */
692 case SpvImageFormatR16
: return 0x822A; /* GL_R16 */
693 case SpvImageFormatR8
: return 0x8229; /* GL_R8 */
694 case SpvImageFormatRgba16Snorm
: return 0x8F9B; /* GL_RGBA16_SNORM */
695 case SpvImageFormatRg16Snorm
: return 0x8F99; /* GL_RG16_SNORM */
696 case SpvImageFormatRg8Snorm
: return 0x8F95; /* GL_RG8_SNORM */
697 case SpvImageFormatR16Snorm
: return 0x8F98; /* GL_R16_SNORM */
698 case SpvImageFormatR8Snorm
: return 0x8F94; /* GL_R8_SNORM */
699 case SpvImageFormatRgba32i
: return 0x8D82; /* GL_RGBA32I */
700 case SpvImageFormatRgba16i
: return 0x8D88; /* GL_RGBA16I */
701 case SpvImageFormatRgba8i
: return 0x8D8E; /* GL_RGBA8I */
702 case SpvImageFormatR32i
: return 0x8235; /* GL_R32I */
703 case SpvImageFormatRg32i
: return 0x823B; /* GL_RG32I */
704 case SpvImageFormatRg16i
: return 0x8239; /* GL_RG16I */
705 case SpvImageFormatRg8i
: return 0x8237; /* GL_RG8I */
706 case SpvImageFormatR16i
: return 0x8233; /* GL_R16I */
707 case SpvImageFormatR8i
: return 0x8231; /* GL_R8I */
708 case SpvImageFormatRgba32ui
: return 0x8D70; /* GL_RGBA32UI */
709 case SpvImageFormatRgba16ui
: return 0x8D76; /* GL_RGBA16UI */
710 case SpvImageFormatRgba8ui
: return 0x8D7C; /* GL_RGBA8UI */
711 case SpvImageFormatR32ui
: return 0x8236; /* GL_R32UI */
712 case SpvImageFormatRgb10a2ui
: return 0x906F; /* GL_RGB10_A2UI */
713 case SpvImageFormatRg32ui
: return 0x823C; /* GL_RG32UI */
714 case SpvImageFormatRg16ui
: return 0x823A; /* GL_RG16UI */
715 case SpvImageFormatRg8ui
: return 0x8238; /* GL_RG8UI */
716 case SpvImageFormatR16ui
: return 0x823A; /* GL_RG16UI */
717 case SpvImageFormatR8ui
: return 0x8232; /* GL_R8UI */
719 assert(!"Invalid image format");
725 vtn_handle_type(struct vtn_builder
*b
, SpvOp opcode
,
726 const uint32_t *w
, unsigned count
)
728 struct vtn_value
*val
= vtn_push_value(b
, w
[1], vtn_value_type_type
);
730 val
->type
= rzalloc(b
, struct vtn_type
);
731 val
->type
->is_builtin
= false;
732 val
->type
->val
= val
;
736 val
->type
->type
= glsl_void_type();
739 val
->type
->type
= glsl_bool_type();
743 const bool signedness
= w
[3];
745 val
->type
->type
= (signedness
? glsl_int64_t_type() : glsl_uint64_t_type());
747 val
->type
->type
= (signedness
? glsl_int_type() : glsl_uint_type());
750 case SpvOpTypeFloat
: {
752 val
->type
->type
= bit_size
== 64 ? glsl_double_type() : glsl_float_type();
756 case SpvOpTypeVector
: {
757 struct vtn_type
*base
= vtn_value(b
, w
[2], vtn_value_type_type
)->type
;
758 unsigned elems
= w
[3];
760 assert(glsl_type_is_scalar(base
->type
));
761 val
->type
->type
= glsl_vector_type(glsl_get_base_type(base
->type
), elems
);
763 /* Vectors implicitly have sizeof(base_type) stride. For now, this
764 * is always 4 bytes. This will have to change if we want to start
765 * supporting doubles or half-floats.
767 val
->type
->stride
= glsl_get_bit_size(base
->type
) / 8;
768 val
->type
->array_element
= base
;
772 case SpvOpTypeMatrix
: {
773 struct vtn_type
*base
= vtn_value(b
, w
[2], vtn_value_type_type
)->type
;
774 unsigned columns
= w
[3];
776 assert(glsl_type_is_vector(base
->type
));
777 val
->type
->type
= glsl_matrix_type(glsl_get_base_type(base
->type
),
778 glsl_get_vector_elements(base
->type
),
780 assert(!glsl_type_is_error(val
->type
->type
));
781 val
->type
->array_element
= base
;
782 val
->type
->row_major
= false;
783 val
->type
->stride
= 0;
787 case SpvOpTypeRuntimeArray
:
788 case SpvOpTypeArray
: {
789 struct vtn_type
*array_element
=
790 vtn_value(b
, w
[2], vtn_value_type_type
)->type
;
793 if (opcode
== SpvOpTypeRuntimeArray
) {
794 /* A length of 0 is used to denote unsized arrays */
798 vtn_value(b
, w
[3], vtn_value_type_constant
)->constant
->values
[0].u32
[0];
801 val
->type
->type
= glsl_array_type(array_element
->type
, length
);
802 val
->type
->array_element
= array_element
;
803 val
->type
->stride
= 0;
807 case SpvOpTypeStruct
: {
808 unsigned num_fields
= count
- 2;
809 val
->type
->members
= ralloc_array(b
, struct vtn_type
*, num_fields
);
810 val
->type
->offsets
= ralloc_array(b
, unsigned, num_fields
);
812 NIR_VLA(struct glsl_struct_field
, fields
, count
);
813 for (unsigned i
= 0; i
< num_fields
; i
++) {
814 val
->type
->members
[i
] =
815 vtn_value(b
, w
[i
+ 2], vtn_value_type_type
)->type
;
816 fields
[i
] = (struct glsl_struct_field
) {
817 .type
= val
->type
->members
[i
]->type
,
818 .name
= ralloc_asprintf(b
, "field%d", i
),
823 struct member_decoration_ctx ctx
= {
824 .num_fields
= num_fields
,
829 vtn_foreach_decoration(b
, val
, struct_member_decoration_cb
, &ctx
);
830 vtn_foreach_decoration(b
, val
, struct_member_matrix_stride_cb
, &ctx
);
832 const char *name
= val
->name
? val
->name
: "struct";
834 val
->type
->type
= glsl_struct_type(fields
, num_fields
, name
);
838 case SpvOpTypeFunction
: {
839 const struct glsl_type
*return_type
=
840 vtn_value(b
, w
[2], vtn_value_type_type
)->type
->type
;
841 NIR_VLA(struct glsl_function_param
, params
, count
- 3);
842 for (unsigned i
= 0; i
< count
- 3; i
++) {
843 params
[i
].type
= vtn_value(b
, w
[i
+ 3], vtn_value_type_type
)->type
->type
;
847 params
[i
].out
= true;
849 val
->type
->type
= glsl_function_type(return_type
, params
, count
- 3);
853 case SpvOpTypePointer
:
854 /* FIXME: For now, we'll just do the really lame thing and return
855 * the same type. The validator should ensure that the proper number
856 * of dereferences happen
858 val
->type
= vtn_value(b
, w
[3], vtn_value_type_type
)->type
;
861 case SpvOpTypeImage
: {
862 const struct glsl_type
*sampled_type
=
863 vtn_value(b
, w
[2], vtn_value_type_type
)->type
->type
;
865 assert(glsl_type_is_vector_or_scalar(sampled_type
));
867 enum glsl_sampler_dim dim
;
868 switch ((SpvDim
)w
[3]) {
869 case SpvDim1D
: dim
= GLSL_SAMPLER_DIM_1D
; break;
870 case SpvDim2D
: dim
= GLSL_SAMPLER_DIM_2D
; break;
871 case SpvDim3D
: dim
= GLSL_SAMPLER_DIM_3D
; break;
872 case SpvDimCube
: dim
= GLSL_SAMPLER_DIM_CUBE
; break;
873 case SpvDimRect
: dim
= GLSL_SAMPLER_DIM_RECT
; break;
874 case SpvDimBuffer
: dim
= GLSL_SAMPLER_DIM_BUF
; break;
875 case SpvDimSubpassData
: dim
= GLSL_SAMPLER_DIM_SUBPASS
; break;
877 unreachable("Invalid SPIR-V Sampler dimension");
880 bool is_shadow
= w
[4];
881 bool is_array
= w
[5];
882 bool multisampled
= w
[6];
883 unsigned sampled
= w
[7];
884 SpvImageFormat format
= w
[8];
887 val
->type
->access_qualifier
= w
[9];
889 val
->type
->access_qualifier
= SpvAccessQualifierReadWrite
;
892 if (dim
== GLSL_SAMPLER_DIM_2D
)
893 dim
= GLSL_SAMPLER_DIM_MS
;
894 else if (dim
== GLSL_SAMPLER_DIM_SUBPASS
)
895 dim
= GLSL_SAMPLER_DIM_SUBPASS_MS
;
897 assert(!"Unsupported multisampled image type");
900 val
->type
->image_format
= translate_image_format(format
);
903 val
->type
->type
= glsl_sampler_type(dim
, is_shadow
, is_array
,
904 glsl_get_base_type(sampled_type
));
905 } else if (sampled
== 2) {
907 val
->type
->type
= glsl_image_type(dim
, is_array
,
908 glsl_get_base_type(sampled_type
));
910 assert(!"We need to know if the image will be sampled");
915 case SpvOpTypeSampledImage
:
916 val
->type
= vtn_value(b
, w
[2], vtn_value_type_type
)->type
;
919 case SpvOpTypeSampler
:
920 /* The actual sampler type here doesn't really matter. It gets
921 * thrown away the moment you combine it with an image. What really
922 * matters is that it's a sampler type as opposed to an integer type
923 * so the backend knows what to do.
925 val
->type
->type
= glsl_bare_sampler_type();
928 case SpvOpTypeOpaque
:
930 case SpvOpTypeDeviceEvent
:
931 case SpvOpTypeReserveId
:
935 unreachable("Unhandled opcode");
938 vtn_foreach_decoration(b
, val
, type_decoration_cb
, NULL
);
941 static nir_constant
*
942 vtn_null_constant(struct vtn_builder
*b
, const struct glsl_type
*type
)
944 nir_constant
*c
= rzalloc(b
, nir_constant
);
946 switch (glsl_get_base_type(type
)) {
949 case GLSL_TYPE_INT64
:
950 case GLSL_TYPE_UINT64
:
952 case GLSL_TYPE_FLOAT
:
953 case GLSL_TYPE_DOUBLE
:
954 /* Nothing to do here. It's already initialized to zero */
957 case GLSL_TYPE_ARRAY
:
958 assert(glsl_get_length(type
) > 0);
959 c
->num_elements
= glsl_get_length(type
);
960 c
->elements
= ralloc_array(b
, nir_constant
*, c
->num_elements
);
962 c
->elements
[0] = vtn_null_constant(b
, glsl_get_array_element(type
));
963 for (unsigned i
= 1; i
< c
->num_elements
; i
++)
964 c
->elements
[i
] = c
->elements
[0];
967 case GLSL_TYPE_STRUCT
:
968 c
->num_elements
= glsl_get_length(type
);
969 c
->elements
= ralloc_array(b
, nir_constant
*, c
->num_elements
);
971 for (unsigned i
= 0; i
< c
->num_elements
; i
++) {
972 c
->elements
[i
] = vtn_null_constant(b
, glsl_get_struct_field(type
, i
));
977 unreachable("Invalid type for null constant");
984 spec_constant_decoration_cb(struct vtn_builder
*b
, struct vtn_value
*v
,
985 int member
, const struct vtn_decoration
*dec
,
988 assert(member
== -1);
989 if (dec
->decoration
!= SpvDecorationSpecId
)
992 struct spec_constant_value
*const_value
= data
;
994 for (unsigned i
= 0; i
< b
->num_specializations
; i
++) {
995 if (b
->specializations
[i
].id
== dec
->literals
[0]) {
996 if (const_value
->is_double
)
997 const_value
->data64
= b
->specializations
[i
].data64
;
999 const_value
->data32
= b
->specializations
[i
].data32
;
1006 get_specialization(struct vtn_builder
*b
, struct vtn_value
*val
,
1007 uint32_t const_value
)
1009 struct spec_constant_value data
;
1010 data
.is_double
= false;
1011 data
.data32
= const_value
;
1012 vtn_foreach_decoration(b
, val
, spec_constant_decoration_cb
, &data
);
1017 get_specialization64(struct vtn_builder
*b
, struct vtn_value
*val
,
1018 uint64_t const_value
)
1020 struct spec_constant_value data
;
1021 data
.is_double
= true;
1022 data
.data64
= const_value
;
1023 vtn_foreach_decoration(b
, val
, spec_constant_decoration_cb
, &data
);
1028 handle_workgroup_size_decoration_cb(struct vtn_builder
*b
,
1029 struct vtn_value
*val
,
1031 const struct vtn_decoration
*dec
,
1034 assert(member
== -1);
1035 if (dec
->decoration
!= SpvDecorationBuiltIn
||
1036 dec
->literals
[0] != SpvBuiltInWorkgroupSize
)
1039 assert(val
->const_type
== glsl_vector_type(GLSL_TYPE_UINT
, 3));
1041 b
->shader
->info
.cs
.local_size
[0] = val
->constant
->values
[0].u32
[0];
1042 b
->shader
->info
.cs
.local_size
[1] = val
->constant
->values
[0].u32
[1];
1043 b
->shader
->info
.cs
.local_size
[2] = val
->constant
->values
[0].u32
[2];
1047 vtn_handle_constant(struct vtn_builder
*b
, SpvOp opcode
,
1048 const uint32_t *w
, unsigned count
)
1050 struct vtn_value
*val
= vtn_push_value(b
, w
[2], vtn_value_type_constant
);
1051 val
->const_type
= vtn_value(b
, w
[1], vtn_value_type_type
)->type
->type
;
1052 val
->constant
= rzalloc(b
, nir_constant
);
1054 case SpvOpConstantTrue
:
1055 assert(val
->const_type
== glsl_bool_type());
1056 val
->constant
->values
[0].u32
[0] = NIR_TRUE
;
1058 case SpvOpConstantFalse
:
1059 assert(val
->const_type
== glsl_bool_type());
1060 val
->constant
->values
[0].u32
[0] = NIR_FALSE
;
1063 case SpvOpSpecConstantTrue
:
1064 case SpvOpSpecConstantFalse
: {
1065 assert(val
->const_type
== glsl_bool_type());
1067 get_specialization(b
, val
, (opcode
== SpvOpSpecConstantTrue
));
1068 val
->constant
->values
[0].u32
[0] = int_val
? NIR_TRUE
: NIR_FALSE
;
1072 case SpvOpConstant
: {
1073 assert(glsl_type_is_scalar(val
->const_type
));
1074 int bit_size
= glsl_get_bit_size(val
->const_type
);
1075 if (bit_size
== 64) {
1076 val
->constant
->values
->u32
[0] = w
[3];
1077 val
->constant
->values
->u32
[1] = w
[4];
1079 assert(bit_size
== 32);
1080 val
->constant
->values
->u32
[0] = w
[3];
1084 case SpvOpSpecConstant
: {
1085 assert(glsl_type_is_scalar(val
->const_type
));
1086 val
->constant
->values
[0].u32
[0] = get_specialization(b
, val
, w
[3]);
1087 int bit_size
= glsl_get_bit_size(val
->const_type
);
1089 val
->constant
->values
[0].u64
[0] =
1090 get_specialization64(b
, val
, vtn_u64_literal(&w
[3]));
1092 val
->constant
->values
[0].u32
[0] = get_specialization(b
, val
, w
[3]);
1095 case SpvOpSpecConstantComposite
:
1096 case SpvOpConstantComposite
: {
1097 unsigned elem_count
= count
- 3;
1098 nir_constant
**elems
= ralloc_array(b
, nir_constant
*, elem_count
);
1099 for (unsigned i
= 0; i
< elem_count
; i
++)
1100 elems
[i
] = vtn_value(b
, w
[i
+ 3], vtn_value_type_constant
)->constant
;
1102 switch (glsl_get_base_type(val
->const_type
)) {
1103 case GLSL_TYPE_UINT
:
1105 case GLSL_TYPE_UINT64
:
1106 case GLSL_TYPE_INT64
:
1107 case GLSL_TYPE_FLOAT
:
1108 case GLSL_TYPE_BOOL
:
1109 case GLSL_TYPE_DOUBLE
: {
1110 int bit_size
= glsl_get_bit_size(val
->const_type
);
1111 if (glsl_type_is_matrix(val
->const_type
)) {
1112 assert(glsl_get_matrix_columns(val
->const_type
) == elem_count
);
1113 for (unsigned i
= 0; i
< elem_count
; i
++)
1114 val
->constant
->values
[i
] = elems
[i
]->values
[0];
1116 assert(glsl_type_is_vector(val
->const_type
));
1117 assert(glsl_get_vector_elements(val
->const_type
) == elem_count
);
1118 for (unsigned i
= 0; i
< elem_count
; i
++) {
1119 if (bit_size
== 64) {
1120 val
->constant
->values
[0].u64
[i
] = elems
[i
]->values
[0].u64
[0];
1122 assert(bit_size
== 32);
1123 val
->constant
->values
[0].u32
[i
] = elems
[i
]->values
[0].u32
[0];
1130 case GLSL_TYPE_STRUCT
:
1131 case GLSL_TYPE_ARRAY
:
1132 ralloc_steal(val
->constant
, elems
);
1133 val
->constant
->num_elements
= elem_count
;
1134 val
->constant
->elements
= elems
;
1138 unreachable("Unsupported type for constants");
1143 case SpvOpSpecConstantOp
: {
1144 SpvOp opcode
= get_specialization(b
, val
, w
[3]);
1146 case SpvOpVectorShuffle
: {
1147 struct vtn_value
*v0
= &b
->values
[w
[4]];
1148 struct vtn_value
*v1
= &b
->values
[w
[5]];
1150 assert(v0
->value_type
== vtn_value_type_constant
||
1151 v0
->value_type
== vtn_value_type_undef
);
1152 assert(v1
->value_type
== vtn_value_type_constant
||
1153 v1
->value_type
== vtn_value_type_undef
);
1155 unsigned len0
= v0
->value_type
== vtn_value_type_constant
?
1156 glsl_get_vector_elements(v0
->const_type
) :
1157 glsl_get_vector_elements(v0
->type
->type
);
1158 unsigned len1
= v1
->value_type
== vtn_value_type_constant
?
1159 glsl_get_vector_elements(v1
->const_type
) :
1160 glsl_get_vector_elements(v1
->type
->type
);
1162 assert(len0
+ len1
< 16);
1164 unsigned bit_size
= glsl_get_bit_size(val
->const_type
);
1165 unsigned bit_size0
= v0
->value_type
== vtn_value_type_constant
?
1166 glsl_get_bit_size(v0
->const_type
) :
1167 glsl_get_bit_size(v0
->type
->type
);
1168 unsigned bit_size1
= v1
->value_type
== vtn_value_type_constant
?
1169 glsl_get_bit_size(v1
->const_type
) :
1170 glsl_get_bit_size(v1
->type
->type
);
1172 assert(bit_size
== bit_size0
&& bit_size
== bit_size1
);
1173 (void)bit_size0
; (void)bit_size1
;
1175 if (bit_size
== 64) {
1177 if (v0
->value_type
== vtn_value_type_constant
) {
1178 for (unsigned i
= 0; i
< len0
; i
++)
1179 u64
[i
] = v0
->constant
->values
[0].u64
[i
];
1181 if (v1
->value_type
== vtn_value_type_constant
) {
1182 for (unsigned i
= 0; i
< len1
; i
++)
1183 u64
[len0
+ i
] = v1
->constant
->values
[0].u64
[i
];
1186 for (unsigned i
= 0, j
= 0; i
< count
- 6; i
++, j
++) {
1187 uint32_t comp
= w
[i
+ 6];
1188 /* If component is not used, set the value to a known constant
1189 * to detect if it is wrongly used.
1191 if (comp
== (uint32_t)-1)
1192 val
->constant
->values
[0].u64
[j
] = 0xdeadbeefdeadbeef;
1194 val
->constant
->values
[0].u64
[j
] = u64
[comp
];
1198 if (v0
->value_type
== vtn_value_type_constant
) {
1199 for (unsigned i
= 0; i
< len0
; i
++)
1200 u32
[i
] = v0
->constant
->values
[0].u32
[i
];
1202 if (v1
->value_type
== vtn_value_type_constant
) {
1203 for (unsigned i
= 0; i
< len1
; i
++)
1204 u32
[len0
+ i
] = v1
->constant
->values
[0].u32
[i
];
1207 for (unsigned i
= 0, j
= 0; i
< count
- 6; i
++, j
++) {
1208 uint32_t comp
= w
[i
+ 6];
1209 /* If component is not used, set the value to a known constant
1210 * to detect if it is wrongly used.
1212 if (comp
== (uint32_t)-1)
1213 val
->constant
->values
[0].u32
[j
] = 0xdeadbeef;
1215 val
->constant
->values
[0].u32
[j
] = u32
[comp
];
1221 case SpvOpCompositeExtract
:
1222 case SpvOpCompositeInsert
: {
1223 struct vtn_value
*comp
;
1224 unsigned deref_start
;
1225 struct nir_constant
**c
;
1226 if (opcode
== SpvOpCompositeExtract
) {
1227 comp
= vtn_value(b
, w
[4], vtn_value_type_constant
);
1229 c
= &comp
->constant
;
1231 comp
= vtn_value(b
, w
[5], vtn_value_type_constant
);
1233 val
->constant
= nir_constant_clone(comp
->constant
,
1240 const struct glsl_type
*type
= comp
->const_type
;
1241 for (unsigned i
= deref_start
; i
< count
; i
++) {
1242 switch (glsl_get_base_type(type
)) {
1243 case GLSL_TYPE_UINT
:
1245 case GLSL_TYPE_UINT64
:
1246 case GLSL_TYPE_INT64
:
1247 case GLSL_TYPE_FLOAT
:
1248 case GLSL_TYPE_DOUBLE
:
1249 case GLSL_TYPE_BOOL
:
1250 /* If we hit this granularity, we're picking off an element */
1251 if (glsl_type_is_matrix(type
)) {
1252 assert(col
== 0 && elem
== -1);
1255 type
= glsl_get_column_type(type
);
1257 assert(elem
<= 0 && glsl_type_is_vector(type
));
1259 type
= glsl_scalar_type(glsl_get_base_type(type
));
1263 case GLSL_TYPE_ARRAY
:
1264 c
= &(*c
)->elements
[w
[i
]];
1265 type
= glsl_get_array_element(type
);
1268 case GLSL_TYPE_STRUCT
:
1269 c
= &(*c
)->elements
[w
[i
]];
1270 type
= glsl_get_struct_field(type
, w
[i
]);
1274 unreachable("Invalid constant type");
1278 if (opcode
== SpvOpCompositeExtract
) {
1282 unsigned num_components
= glsl_get_vector_elements(type
);
1283 unsigned bit_size
= glsl_get_bit_size(type
);
1284 for (unsigned i
= 0; i
< num_components
; i
++)
1285 if (bit_size
== 64) {
1286 val
->constant
->values
[0].u64
[i
] = (*c
)->values
[col
].u64
[elem
+ i
];
1288 assert(bit_size
== 32);
1289 val
->constant
->values
[0].u32
[i
] = (*c
)->values
[col
].u32
[elem
+ i
];
1293 struct vtn_value
*insert
=
1294 vtn_value(b
, w
[4], vtn_value_type_constant
);
1295 assert(insert
->const_type
== type
);
1297 *c
= insert
->constant
;
1299 unsigned num_components
= glsl_get_vector_elements(type
);
1300 unsigned bit_size
= glsl_get_bit_size(type
);
1301 for (unsigned i
= 0; i
< num_components
; i
++)
1302 if (bit_size
== 64) {
1303 (*c
)->values
[col
].u64
[elem
+ i
] = insert
->constant
->values
[0].u64
[i
];
1305 assert(bit_size
== 32);
1306 (*c
)->values
[col
].u32
[elem
+ i
] = insert
->constant
->values
[0].u32
[i
];
1315 nir_alu_type dst_alu_type
= nir_get_nir_type_for_glsl_type(val
->const_type
);
1316 nir_alu_type src_alu_type
= dst_alu_type
;
1317 nir_op op
= vtn_nir_alu_op_for_spirv_opcode(opcode
, &swap
, src_alu_type
, dst_alu_type
);
1319 unsigned num_components
= glsl_get_vector_elements(val
->const_type
);
1321 glsl_get_bit_size(val
->const_type
);
1323 nir_const_value src
[4];
1325 for (unsigned i
= 0; i
< count
- 4; i
++) {
1327 vtn_value(b
, w
[4 + i
], vtn_value_type_constant
)->constant
;
1329 unsigned j
= swap
? 1 - i
: i
;
1330 assert(bit_size
== 32);
1331 src
[j
] = c
->values
[0];
1334 val
->constant
->values
[0] =
1335 nir_eval_const_opcode(op
, num_components
, bit_size
, src
);
1342 case SpvOpConstantNull
:
1343 val
->constant
= vtn_null_constant(b
, val
->const_type
);
1346 case SpvOpConstantSampler
:
1347 assert(!"OpConstantSampler requires Kernel Capability");
1351 unreachable("Unhandled opcode");
1354 /* Now that we have the value, update the workgroup size if needed */
1355 vtn_foreach_decoration(b
, val
, handle_workgroup_size_decoration_cb
, NULL
);
1359 vtn_handle_function_call(struct vtn_builder
*b
, SpvOp opcode
,
1360 const uint32_t *w
, unsigned count
)
1362 struct nir_function
*callee
=
1363 vtn_value(b
, w
[3], vtn_value_type_function
)->func
->impl
->function
;
1365 nir_call_instr
*call
= nir_call_instr_create(b
->nb
.shader
, callee
);
1366 for (unsigned i
= 0; i
< call
->num_params
; i
++) {
1367 unsigned arg_id
= w
[4 + i
];
1368 struct vtn_value
*arg
= vtn_untyped_value(b
, arg_id
);
1369 if (arg
->value_type
== vtn_value_type_pointer
) {
1370 nir_deref_var
*d
= vtn_pointer_to_deref(b
, arg
->pointer
);
1371 call
->params
[i
] = nir_deref_var_clone(d
, call
);
1373 struct vtn_ssa_value
*arg_ssa
= vtn_ssa_value(b
, arg_id
);
1375 /* Make a temporary to store the argument in */
1377 nir_local_variable_create(b
->impl
, arg_ssa
->type
, "arg_tmp");
1378 call
->params
[i
] = nir_deref_var_create(call
, tmp
);
1380 vtn_local_store(b
, arg_ssa
, call
->params
[i
]);
1384 nir_variable
*out_tmp
= NULL
;
1385 if (!glsl_type_is_void(callee
->return_type
)) {
1386 out_tmp
= nir_local_variable_create(b
->impl
, callee
->return_type
,
1388 call
->return_deref
= nir_deref_var_create(call
, out_tmp
);
1391 nir_builder_instr_insert(&b
->nb
, &call
->instr
);
1393 if (glsl_type_is_void(callee
->return_type
)) {
1394 vtn_push_value(b
, w
[2], vtn_value_type_undef
);
1396 struct vtn_value
*retval
= vtn_push_value(b
, w
[2], vtn_value_type_ssa
);
1397 retval
->ssa
= vtn_local_load(b
, call
->return_deref
);
1401 struct vtn_ssa_value
*
1402 vtn_create_ssa_value(struct vtn_builder
*b
, const struct glsl_type
*type
)
1404 struct vtn_ssa_value
*val
= rzalloc(b
, struct vtn_ssa_value
);
1407 if (!glsl_type_is_vector_or_scalar(type
)) {
1408 unsigned elems
= glsl_get_length(type
);
1409 val
->elems
= ralloc_array(b
, struct vtn_ssa_value
*, elems
);
1410 for (unsigned i
= 0; i
< elems
; i
++) {
1411 const struct glsl_type
*child_type
;
1413 switch (glsl_get_base_type(type
)) {
1415 case GLSL_TYPE_UINT
:
1416 case GLSL_TYPE_INT64
:
1417 case GLSL_TYPE_UINT64
:
1418 case GLSL_TYPE_BOOL
:
1419 case GLSL_TYPE_FLOAT
:
1420 case GLSL_TYPE_DOUBLE
:
1421 child_type
= glsl_get_column_type(type
);
1423 case GLSL_TYPE_ARRAY
:
1424 child_type
= glsl_get_array_element(type
);
1426 case GLSL_TYPE_STRUCT
:
1427 child_type
= glsl_get_struct_field(type
, i
);
1430 unreachable("unkown base type");
1433 val
->elems
[i
] = vtn_create_ssa_value(b
, child_type
);
1441 vtn_tex_src(struct vtn_builder
*b
, unsigned index
, nir_tex_src_type type
)
1444 src
.src
= nir_src_for_ssa(vtn_ssa_value(b
, index
)->def
);
1445 src
.src_type
= type
;
1450 vtn_handle_texture(struct vtn_builder
*b
, SpvOp opcode
,
1451 const uint32_t *w
, unsigned count
)
1453 if (opcode
== SpvOpSampledImage
) {
1454 struct vtn_value
*val
=
1455 vtn_push_value(b
, w
[2], vtn_value_type_sampled_image
);
1456 val
->sampled_image
= ralloc(b
, struct vtn_sampled_image
);
1457 val
->sampled_image
->image
=
1458 vtn_value(b
, w
[3], vtn_value_type_pointer
)->pointer
;
1459 val
->sampled_image
->sampler
=
1460 vtn_value(b
, w
[4], vtn_value_type_pointer
)->pointer
;
1462 } else if (opcode
== SpvOpImage
) {
1463 struct vtn_value
*val
= vtn_push_value(b
, w
[2], vtn_value_type_pointer
);
1464 struct vtn_value
*src_val
= vtn_untyped_value(b
, w
[3]);
1465 if (src_val
->value_type
== vtn_value_type_sampled_image
) {
1466 val
->pointer
= src_val
->sampled_image
->image
;
1468 assert(src_val
->value_type
== vtn_value_type_pointer
);
1469 val
->pointer
= src_val
->pointer
;
1474 struct vtn_type
*ret_type
= vtn_value(b
, w
[1], vtn_value_type_type
)->type
;
1475 struct vtn_value
*val
= vtn_push_value(b
, w
[2], vtn_value_type_ssa
);
1477 struct vtn_sampled_image sampled
;
1478 struct vtn_value
*sampled_val
= vtn_untyped_value(b
, w
[3]);
1479 if (sampled_val
->value_type
== vtn_value_type_sampled_image
) {
1480 sampled
= *sampled_val
->sampled_image
;
1482 assert(sampled_val
->value_type
== vtn_value_type_pointer
);
1483 sampled
.image
= NULL
;
1484 sampled
.sampler
= sampled_val
->pointer
;
1487 const struct glsl_type
*image_type
;
1488 if (sampled
.image
) {
1489 image_type
= sampled
.image
->var
->var
->interface_type
;
1491 image_type
= sampled
.sampler
->var
->var
->interface_type
;
1493 const enum glsl_sampler_dim sampler_dim
= glsl_get_sampler_dim(image_type
);
1494 const bool is_array
= glsl_sampler_type_is_array(image_type
);
1495 const bool is_shadow
= glsl_sampler_type_is_shadow(image_type
);
1497 /* Figure out the base texture operation */
1500 case SpvOpImageSampleImplicitLod
:
1501 case SpvOpImageSampleDrefImplicitLod
:
1502 case SpvOpImageSampleProjImplicitLod
:
1503 case SpvOpImageSampleProjDrefImplicitLod
:
1504 texop
= nir_texop_tex
;
1507 case SpvOpImageSampleExplicitLod
:
1508 case SpvOpImageSampleDrefExplicitLod
:
1509 case SpvOpImageSampleProjExplicitLod
:
1510 case SpvOpImageSampleProjDrefExplicitLod
:
1511 texop
= nir_texop_txl
;
1514 case SpvOpImageFetch
:
1515 if (glsl_get_sampler_dim(image_type
) == GLSL_SAMPLER_DIM_MS
) {
1516 texop
= nir_texop_txf_ms
;
1518 texop
= nir_texop_txf
;
1522 case SpvOpImageGather
:
1523 case SpvOpImageDrefGather
:
1524 texop
= nir_texop_tg4
;
1527 case SpvOpImageQuerySizeLod
:
1528 case SpvOpImageQuerySize
:
1529 texop
= nir_texop_txs
;
1532 case SpvOpImageQueryLod
:
1533 texop
= nir_texop_lod
;
1536 case SpvOpImageQueryLevels
:
1537 texop
= nir_texop_query_levels
;
1540 case SpvOpImageQuerySamples
:
1541 texop
= nir_texop_texture_samples
;
1545 unreachable("Unhandled opcode");
1548 nir_tex_src srcs
[8]; /* 8 should be enough */
1549 nir_tex_src
*p
= srcs
;
1553 struct nir_ssa_def
*coord
;
1554 unsigned coord_components
;
1556 case SpvOpImageSampleImplicitLod
:
1557 case SpvOpImageSampleExplicitLod
:
1558 case SpvOpImageSampleDrefImplicitLod
:
1559 case SpvOpImageSampleDrefExplicitLod
:
1560 case SpvOpImageSampleProjImplicitLod
:
1561 case SpvOpImageSampleProjExplicitLod
:
1562 case SpvOpImageSampleProjDrefImplicitLod
:
1563 case SpvOpImageSampleProjDrefExplicitLod
:
1564 case SpvOpImageFetch
:
1565 case SpvOpImageGather
:
1566 case SpvOpImageDrefGather
:
1567 case SpvOpImageQueryLod
: {
1568 /* All these types have the coordinate as their first real argument */
1569 switch (sampler_dim
) {
1570 case GLSL_SAMPLER_DIM_1D
:
1571 case GLSL_SAMPLER_DIM_BUF
:
1572 coord_components
= 1;
1574 case GLSL_SAMPLER_DIM_2D
:
1575 case GLSL_SAMPLER_DIM_RECT
:
1576 case GLSL_SAMPLER_DIM_MS
:
1577 coord_components
= 2;
1579 case GLSL_SAMPLER_DIM_3D
:
1580 case GLSL_SAMPLER_DIM_CUBE
:
1581 coord_components
= 3;
1584 unreachable("Invalid sampler type");
1587 if (is_array
&& texop
!= nir_texop_lod
)
1590 coord
= vtn_ssa_value(b
, w
[idx
++])->def
;
1591 p
->src
= nir_src_for_ssa(nir_channels(&b
->nb
, coord
,
1592 (1 << coord_components
) - 1));
1593 p
->src_type
= nir_tex_src_coord
;
1600 coord_components
= 0;
1605 case SpvOpImageSampleProjImplicitLod
:
1606 case SpvOpImageSampleProjExplicitLod
:
1607 case SpvOpImageSampleProjDrefImplicitLod
:
1608 case SpvOpImageSampleProjDrefExplicitLod
:
1609 /* These have the projector as the last coordinate component */
1610 p
->src
= nir_src_for_ssa(nir_channel(&b
->nb
, coord
, coord_components
));
1611 p
->src_type
= nir_tex_src_projector
;
1619 unsigned gather_component
= 0;
1621 case SpvOpImageSampleDrefImplicitLod
:
1622 case SpvOpImageSampleDrefExplicitLod
:
1623 case SpvOpImageSampleProjDrefImplicitLod
:
1624 case SpvOpImageSampleProjDrefExplicitLod
:
1625 case SpvOpImageDrefGather
:
1626 /* These all have an explicit depth value as their next source */
1627 (*p
++) = vtn_tex_src(b
, w
[idx
++], nir_tex_src_comparator
);
1630 case SpvOpImageGather
:
1631 /* This has a component as its next source */
1633 vtn_value(b
, w
[idx
++], vtn_value_type_constant
)->constant
->values
[0].u32
[0];
1640 /* For OpImageQuerySizeLod, we always have an LOD */
1641 if (opcode
== SpvOpImageQuerySizeLod
)
1642 (*p
++) = vtn_tex_src(b
, w
[idx
++], nir_tex_src_lod
);
1644 /* Now we need to handle some number of optional arguments */
1645 const struct vtn_ssa_value
*gather_offsets
= NULL
;
1647 uint32_t operands
= w
[idx
++];
1649 if (operands
& SpvImageOperandsBiasMask
) {
1650 assert(texop
== nir_texop_tex
);
1651 texop
= nir_texop_txb
;
1652 (*p
++) = vtn_tex_src(b
, w
[idx
++], nir_tex_src_bias
);
1655 if (operands
& SpvImageOperandsLodMask
) {
1656 assert(texop
== nir_texop_txl
|| texop
== nir_texop_txf
||
1657 texop
== nir_texop_txs
);
1658 (*p
++) = vtn_tex_src(b
, w
[idx
++], nir_tex_src_lod
);
1661 if (operands
& SpvImageOperandsGradMask
) {
1662 assert(texop
== nir_texop_txl
);
1663 texop
= nir_texop_txd
;
1664 (*p
++) = vtn_tex_src(b
, w
[idx
++], nir_tex_src_ddx
);
1665 (*p
++) = vtn_tex_src(b
, w
[idx
++], nir_tex_src_ddy
);
1668 if (operands
& SpvImageOperandsOffsetMask
||
1669 operands
& SpvImageOperandsConstOffsetMask
)
1670 (*p
++) = vtn_tex_src(b
, w
[idx
++], nir_tex_src_offset
);
1672 if (operands
& SpvImageOperandsConstOffsetsMask
) {
1673 gather_offsets
= vtn_ssa_value(b
, w
[idx
++]);
1674 (*p
++) = (nir_tex_src
){};
1677 if (operands
& SpvImageOperandsSampleMask
) {
1678 assert(texop
== nir_texop_txf_ms
);
1679 texop
= nir_texop_txf_ms
;
1680 (*p
++) = vtn_tex_src(b
, w
[idx
++], nir_tex_src_ms_index
);
1683 /* We should have now consumed exactly all of the arguments */
1684 assert(idx
== count
);
1686 nir_tex_instr
*instr
= nir_tex_instr_create(b
->shader
, p
- srcs
);
1689 memcpy(instr
->src
, srcs
, instr
->num_srcs
* sizeof(*instr
->src
));
1691 instr
->coord_components
= coord_components
;
1692 instr
->sampler_dim
= sampler_dim
;
1693 instr
->is_array
= is_array
;
1694 instr
->is_shadow
= is_shadow
;
1695 instr
->is_new_style_shadow
=
1696 is_shadow
&& glsl_get_components(ret_type
->type
) == 1;
1697 instr
->component
= gather_component
;
1699 switch (glsl_get_sampler_result_type(image_type
)) {
1700 case GLSL_TYPE_FLOAT
: instr
->dest_type
= nir_type_float
; break;
1701 case GLSL_TYPE_INT
: instr
->dest_type
= nir_type_int
; break;
1702 case GLSL_TYPE_UINT
: instr
->dest_type
= nir_type_uint
; break;
1703 case GLSL_TYPE_BOOL
: instr
->dest_type
= nir_type_bool
; break;
1705 unreachable("Invalid base type for sampler result");
1708 nir_deref_var
*sampler
= vtn_pointer_to_deref(b
, sampled
.sampler
);
1709 nir_deref_var
*texture
;
1710 if (sampled
.image
) {
1711 nir_deref_var
*image
= vtn_pointer_to_deref(b
, sampled
.image
);
1717 instr
->texture
= nir_deref_var_clone(texture
, instr
);
1719 switch (instr
->op
) {
1724 /* These operations require a sampler */
1725 instr
->sampler
= nir_deref_var_clone(sampler
, instr
);
1728 case nir_texop_txf_ms
:
1732 case nir_texop_query_levels
:
1733 case nir_texop_texture_samples
:
1734 case nir_texop_samples_identical
:
1736 instr
->sampler
= NULL
;
1738 case nir_texop_txf_ms_mcs
:
1739 unreachable("unexpected nir_texop_txf_ms_mcs");
1742 nir_ssa_dest_init(&instr
->instr
, &instr
->dest
,
1743 nir_tex_instr_dest_size(instr
), 32, NULL
);
1745 assert(glsl_get_vector_elements(ret_type
->type
) ==
1746 nir_tex_instr_dest_size(instr
));
1749 nir_instr
*instruction
;
1750 if (gather_offsets
) {
1751 assert(glsl_get_base_type(gather_offsets
->type
) == GLSL_TYPE_ARRAY
);
1752 assert(glsl_get_length(gather_offsets
->type
) == 4);
1753 nir_tex_instr
*instrs
[4] = {instr
, NULL
, NULL
, NULL
};
1755 /* Copy the current instruction 4x */
1756 for (uint32_t i
= 1; i
< 4; i
++) {
1757 instrs
[i
] = nir_tex_instr_create(b
->shader
, instr
->num_srcs
);
1758 instrs
[i
]->op
= instr
->op
;
1759 instrs
[i
]->coord_components
= instr
->coord_components
;
1760 instrs
[i
]->sampler_dim
= instr
->sampler_dim
;
1761 instrs
[i
]->is_array
= instr
->is_array
;
1762 instrs
[i
]->is_shadow
= instr
->is_shadow
;
1763 instrs
[i
]->is_new_style_shadow
= instr
->is_new_style_shadow
;
1764 instrs
[i
]->component
= instr
->component
;
1765 instrs
[i
]->dest_type
= instr
->dest_type
;
1766 instrs
[i
]->texture
= nir_deref_var_clone(texture
, instrs
[i
]);
1767 instrs
[i
]->sampler
= NULL
;
1769 memcpy(instrs
[i
]->src
, srcs
, instr
->num_srcs
* sizeof(*instr
->src
));
1771 nir_ssa_dest_init(&instrs
[i
]->instr
, &instrs
[i
]->dest
,
1772 nir_tex_instr_dest_size(instr
), 32, NULL
);
1775 /* Fill in the last argument with the offset from the passed in offsets
1776 * and insert the instruction into the stream.
1778 for (uint32_t i
= 0; i
< 4; i
++) {
1780 src
.src
= nir_src_for_ssa(gather_offsets
->elems
[i
]->def
);
1781 src
.src_type
= nir_tex_src_offset
;
1782 instrs
[i
]->src
[instrs
[i
]->num_srcs
- 1] = src
;
1783 nir_builder_instr_insert(&b
->nb
, &instrs
[i
]->instr
);
1786 /* Combine the results of the 4 instructions by taking their .w
1789 nir_alu_instr
*vec4
= nir_alu_instr_create(b
->shader
, nir_op_vec4
);
1790 nir_ssa_dest_init(&vec4
->instr
, &vec4
->dest
.dest
, 4, 32, NULL
);
1791 vec4
->dest
.write_mask
= 0xf;
1792 for (uint32_t i
= 0; i
< 4; i
++) {
1793 vec4
->src
[i
].src
= nir_src_for_ssa(&instrs
[i
]->dest
.ssa
);
1794 vec4
->src
[i
].swizzle
[0] = 3;
1796 def
= &vec4
->dest
.dest
.ssa
;
1797 instruction
= &vec4
->instr
;
1799 def
= &instr
->dest
.ssa
;
1800 instruction
= &instr
->instr
;
1803 val
->ssa
= vtn_create_ssa_value(b
, ret_type
->type
);
1804 val
->ssa
->def
= def
;
1806 nir_builder_instr_insert(&b
->nb
, instruction
);
1810 fill_common_atomic_sources(struct vtn_builder
*b
, SpvOp opcode
,
1811 const uint32_t *w
, nir_src
*src
)
1814 case SpvOpAtomicIIncrement
:
1815 src
[0] = nir_src_for_ssa(nir_imm_int(&b
->nb
, 1));
1818 case SpvOpAtomicIDecrement
:
1819 src
[0] = nir_src_for_ssa(nir_imm_int(&b
->nb
, -1));
1822 case SpvOpAtomicISub
:
1824 nir_src_for_ssa(nir_ineg(&b
->nb
, vtn_ssa_value(b
, w
[6])->def
));
1827 case SpvOpAtomicCompareExchange
:
1828 src
[0] = nir_src_for_ssa(vtn_ssa_value(b
, w
[8])->def
);
1829 src
[1] = nir_src_for_ssa(vtn_ssa_value(b
, w
[7])->def
);
1832 case SpvOpAtomicExchange
:
1833 case SpvOpAtomicIAdd
:
1834 case SpvOpAtomicSMin
:
1835 case SpvOpAtomicUMin
:
1836 case SpvOpAtomicSMax
:
1837 case SpvOpAtomicUMax
:
1838 case SpvOpAtomicAnd
:
1840 case SpvOpAtomicXor
:
1841 src
[0] = nir_src_for_ssa(vtn_ssa_value(b
, w
[6])->def
);
1845 unreachable("Invalid SPIR-V atomic");
1849 static nir_ssa_def
*
1850 get_image_coord(struct vtn_builder
*b
, uint32_t value
)
1852 struct vtn_ssa_value
*coord
= vtn_ssa_value(b
, value
);
1854 /* The image_load_store intrinsics assume a 4-dim coordinate */
1855 unsigned dim
= glsl_get_vector_elements(coord
->type
);
1856 unsigned swizzle
[4];
1857 for (unsigned i
= 0; i
< 4; i
++)
1858 swizzle
[i
] = MIN2(i
, dim
- 1);
1860 return nir_swizzle(&b
->nb
, coord
->def
, swizzle
, 4, false);
1864 vtn_handle_image(struct vtn_builder
*b
, SpvOp opcode
,
1865 const uint32_t *w
, unsigned count
)
1867 /* Just get this one out of the way */
1868 if (opcode
== SpvOpImageTexelPointer
) {
1869 struct vtn_value
*val
=
1870 vtn_push_value(b
, w
[2], vtn_value_type_image_pointer
);
1871 val
->image
= ralloc(b
, struct vtn_image_pointer
);
1873 val
->image
->image
= vtn_value(b
, w
[3], vtn_value_type_pointer
)->pointer
;
1874 val
->image
->coord
= get_image_coord(b
, w
[4]);
1875 val
->image
->sample
= vtn_ssa_value(b
, w
[5])->def
;
1879 struct vtn_image_pointer image
;
1882 case SpvOpAtomicExchange
:
1883 case SpvOpAtomicCompareExchange
:
1884 case SpvOpAtomicCompareExchangeWeak
:
1885 case SpvOpAtomicIIncrement
:
1886 case SpvOpAtomicIDecrement
:
1887 case SpvOpAtomicIAdd
:
1888 case SpvOpAtomicISub
:
1889 case SpvOpAtomicLoad
:
1890 case SpvOpAtomicSMin
:
1891 case SpvOpAtomicUMin
:
1892 case SpvOpAtomicSMax
:
1893 case SpvOpAtomicUMax
:
1894 case SpvOpAtomicAnd
:
1896 case SpvOpAtomicXor
:
1897 image
= *vtn_value(b
, w
[3], vtn_value_type_image_pointer
)->image
;
1900 case SpvOpAtomicStore
:
1901 image
= *vtn_value(b
, w
[1], vtn_value_type_image_pointer
)->image
;
1904 case SpvOpImageQuerySize
:
1905 image
.image
= vtn_value(b
, w
[3], vtn_value_type_pointer
)->pointer
;
1907 image
.sample
= NULL
;
1910 case SpvOpImageRead
:
1911 image
.image
= vtn_value(b
, w
[3], vtn_value_type_pointer
)->pointer
;
1912 image
.coord
= get_image_coord(b
, w
[4]);
1914 if (count
> 5 && (w
[5] & SpvImageOperandsSampleMask
)) {
1915 assert(w
[5] == SpvImageOperandsSampleMask
);
1916 image
.sample
= vtn_ssa_value(b
, w
[6])->def
;
1918 image
.sample
= nir_ssa_undef(&b
->nb
, 1, 32);
1922 case SpvOpImageWrite
:
1923 image
.image
= vtn_value(b
, w
[1], vtn_value_type_pointer
)->pointer
;
1924 image
.coord
= get_image_coord(b
, w
[2]);
1928 if (count
> 4 && (w
[4] & SpvImageOperandsSampleMask
)) {
1929 assert(w
[4] == SpvImageOperandsSampleMask
);
1930 image
.sample
= vtn_ssa_value(b
, w
[5])->def
;
1932 image
.sample
= nir_ssa_undef(&b
->nb
, 1, 32);
1937 unreachable("Invalid image opcode");
1940 nir_intrinsic_op op
;
1942 #define OP(S, N) case SpvOp##S: op = nir_intrinsic_image_##N; break;
1943 OP(ImageQuerySize
, size
)
1945 OP(ImageWrite
, store
)
1946 OP(AtomicLoad
, load
)
1947 OP(AtomicStore
, store
)
1948 OP(AtomicExchange
, atomic_exchange
)
1949 OP(AtomicCompareExchange
, atomic_comp_swap
)
1950 OP(AtomicIIncrement
, atomic_add
)
1951 OP(AtomicIDecrement
, atomic_add
)
1952 OP(AtomicIAdd
, atomic_add
)
1953 OP(AtomicISub
, atomic_add
)
1954 OP(AtomicSMin
, atomic_min
)
1955 OP(AtomicUMin
, atomic_min
)
1956 OP(AtomicSMax
, atomic_max
)
1957 OP(AtomicUMax
, atomic_max
)
1958 OP(AtomicAnd
, atomic_and
)
1959 OP(AtomicOr
, atomic_or
)
1960 OP(AtomicXor
, atomic_xor
)
1963 unreachable("Invalid image opcode");
1966 nir_intrinsic_instr
*intrin
= nir_intrinsic_instr_create(b
->shader
, op
);
1968 nir_deref_var
*image_deref
= vtn_pointer_to_deref(b
, image
.image
);
1969 intrin
->variables
[0] = nir_deref_var_clone(image_deref
, intrin
);
1971 /* ImageQuerySize doesn't take any extra parameters */
1972 if (opcode
!= SpvOpImageQuerySize
) {
1973 /* The image coordinate is always 4 components but we may not have that
1974 * many. Swizzle to compensate.
1977 for (unsigned i
= 0; i
< 4; i
++)
1978 swiz
[i
] = i
< image
.coord
->num_components
? i
: 0;
1979 intrin
->src
[0] = nir_src_for_ssa(nir_swizzle(&b
->nb
, image
.coord
,
1981 intrin
->src
[1] = nir_src_for_ssa(image
.sample
);
1985 case SpvOpAtomicLoad
:
1986 case SpvOpImageQuerySize
:
1987 case SpvOpImageRead
:
1989 case SpvOpAtomicStore
:
1990 intrin
->src
[2] = nir_src_for_ssa(vtn_ssa_value(b
, w
[4])->def
);
1992 case SpvOpImageWrite
:
1993 intrin
->src
[2] = nir_src_for_ssa(vtn_ssa_value(b
, w
[3])->def
);
1996 case SpvOpAtomicIIncrement
:
1997 case SpvOpAtomicIDecrement
:
1998 case SpvOpAtomicExchange
:
1999 case SpvOpAtomicIAdd
:
2000 case SpvOpAtomicSMin
:
2001 case SpvOpAtomicUMin
:
2002 case SpvOpAtomicSMax
:
2003 case SpvOpAtomicUMax
:
2004 case SpvOpAtomicAnd
:
2006 case SpvOpAtomicXor
:
2007 fill_common_atomic_sources(b
, opcode
, w
, &intrin
->src
[2]);
2011 unreachable("Invalid image opcode");
2014 if (opcode
!= SpvOpImageWrite
) {
2015 struct vtn_value
*val
= vtn_push_value(b
, w
[2], vtn_value_type_ssa
);
2016 struct vtn_type
*type
= vtn_value(b
, w
[1], vtn_value_type_type
)->type
;
2018 unsigned dest_components
=
2019 nir_intrinsic_infos
[intrin
->intrinsic
].dest_components
;
2020 if (intrin
->intrinsic
== nir_intrinsic_image_size
) {
2021 dest_components
= intrin
->num_components
=
2022 glsl_get_vector_elements(type
->type
);
2025 nir_ssa_dest_init(&intrin
->instr
, &intrin
->dest
,
2026 dest_components
, 32, NULL
);
2028 nir_builder_instr_insert(&b
->nb
, &intrin
->instr
);
2030 val
->ssa
= vtn_create_ssa_value(b
, type
->type
);
2031 val
->ssa
->def
= &intrin
->dest
.ssa
;
2033 nir_builder_instr_insert(&b
->nb
, &intrin
->instr
);
2037 static nir_intrinsic_op
2038 get_ssbo_nir_atomic_op(SpvOp opcode
)
2041 case SpvOpAtomicLoad
: return nir_intrinsic_load_ssbo
;
2042 case SpvOpAtomicStore
: return nir_intrinsic_store_ssbo
;
2043 #define OP(S, N) case SpvOp##S: return nir_intrinsic_ssbo_##N;
2044 OP(AtomicExchange
, atomic_exchange
)
2045 OP(AtomicCompareExchange
, atomic_comp_swap
)
2046 OP(AtomicIIncrement
, atomic_add
)
2047 OP(AtomicIDecrement
, atomic_add
)
2048 OP(AtomicIAdd
, atomic_add
)
2049 OP(AtomicISub
, atomic_add
)
2050 OP(AtomicSMin
, atomic_imin
)
2051 OP(AtomicUMin
, atomic_umin
)
2052 OP(AtomicSMax
, atomic_imax
)
2053 OP(AtomicUMax
, atomic_umax
)
2054 OP(AtomicAnd
, atomic_and
)
2055 OP(AtomicOr
, atomic_or
)
2056 OP(AtomicXor
, atomic_xor
)
2059 unreachable("Invalid SSBO atomic");
2063 static nir_intrinsic_op
2064 get_shared_nir_atomic_op(SpvOp opcode
)
2067 case SpvOpAtomicLoad
: return nir_intrinsic_load_var
;
2068 case SpvOpAtomicStore
: return nir_intrinsic_store_var
;
2069 #define OP(S, N) case SpvOp##S: return nir_intrinsic_var_##N;
2070 OP(AtomicExchange
, atomic_exchange
)
2071 OP(AtomicCompareExchange
, atomic_comp_swap
)
2072 OP(AtomicIIncrement
, atomic_add
)
2073 OP(AtomicIDecrement
, atomic_add
)
2074 OP(AtomicIAdd
, atomic_add
)
2075 OP(AtomicISub
, atomic_add
)
2076 OP(AtomicSMin
, atomic_imin
)
2077 OP(AtomicUMin
, atomic_umin
)
2078 OP(AtomicSMax
, atomic_imax
)
2079 OP(AtomicUMax
, atomic_umax
)
2080 OP(AtomicAnd
, atomic_and
)
2081 OP(AtomicOr
, atomic_or
)
2082 OP(AtomicXor
, atomic_xor
)
2085 unreachable("Invalid shared atomic");
2090 vtn_handle_ssbo_or_shared_atomic(struct vtn_builder
*b
, SpvOp opcode
,
2091 const uint32_t *w
, unsigned count
)
2093 struct vtn_pointer
*ptr
;
2094 nir_intrinsic_instr
*atomic
;
2097 case SpvOpAtomicLoad
:
2098 case SpvOpAtomicExchange
:
2099 case SpvOpAtomicCompareExchange
:
2100 case SpvOpAtomicCompareExchangeWeak
:
2101 case SpvOpAtomicIIncrement
:
2102 case SpvOpAtomicIDecrement
:
2103 case SpvOpAtomicIAdd
:
2104 case SpvOpAtomicISub
:
2105 case SpvOpAtomicSMin
:
2106 case SpvOpAtomicUMin
:
2107 case SpvOpAtomicSMax
:
2108 case SpvOpAtomicUMax
:
2109 case SpvOpAtomicAnd
:
2111 case SpvOpAtomicXor
:
2112 ptr
= vtn_value(b
, w
[3], vtn_value_type_pointer
)->pointer
;
2115 case SpvOpAtomicStore
:
2116 ptr
= vtn_value(b
, w
[1], vtn_value_type_pointer
)->pointer
;
2120 unreachable("Invalid SPIR-V atomic");
2124 SpvScope scope = w[4];
2125 SpvMemorySemanticsMask semantics = w[5];
2128 if (ptr
->mode
== vtn_variable_mode_workgroup
) {
2129 nir_deref_var
*deref
= vtn_pointer_to_deref(b
, ptr
);
2130 const struct glsl_type
*deref_type
= nir_deref_tail(&deref
->deref
)->type
;
2131 nir_intrinsic_op op
= get_shared_nir_atomic_op(opcode
);
2132 atomic
= nir_intrinsic_instr_create(b
->nb
.shader
, op
);
2133 atomic
->variables
[0] = nir_deref_var_clone(deref
, atomic
);
2136 case SpvOpAtomicLoad
:
2137 atomic
->num_components
= glsl_get_vector_elements(deref_type
);
2140 case SpvOpAtomicStore
:
2141 atomic
->num_components
= glsl_get_vector_elements(deref_type
);
2142 nir_intrinsic_set_write_mask(atomic
, (1 << atomic
->num_components
) - 1);
2143 atomic
->src
[0] = nir_src_for_ssa(vtn_ssa_value(b
, w
[4])->def
);
2146 case SpvOpAtomicExchange
:
2147 case SpvOpAtomicCompareExchange
:
2148 case SpvOpAtomicCompareExchangeWeak
:
2149 case SpvOpAtomicIIncrement
:
2150 case SpvOpAtomicIDecrement
:
2151 case SpvOpAtomicIAdd
:
2152 case SpvOpAtomicISub
:
2153 case SpvOpAtomicSMin
:
2154 case SpvOpAtomicUMin
:
2155 case SpvOpAtomicSMax
:
2156 case SpvOpAtomicUMax
:
2157 case SpvOpAtomicAnd
:
2159 case SpvOpAtomicXor
:
2160 fill_common_atomic_sources(b
, opcode
, w
, &atomic
->src
[0]);
2164 unreachable("Invalid SPIR-V atomic");
2168 assert(ptr
->mode
== vtn_variable_mode_ssbo
);
2169 nir_ssa_def
*offset
, *index
;
2170 offset
= vtn_pointer_to_offset(b
, ptr
, &index
, NULL
);
2172 nir_intrinsic_op op
= get_ssbo_nir_atomic_op(opcode
);
2174 atomic
= nir_intrinsic_instr_create(b
->nb
.shader
, op
);
2177 case SpvOpAtomicLoad
:
2178 atomic
->num_components
= glsl_get_vector_elements(ptr
->type
->type
);
2179 atomic
->src
[0] = nir_src_for_ssa(index
);
2180 atomic
->src
[1] = nir_src_for_ssa(offset
);
2183 case SpvOpAtomicStore
:
2184 atomic
->num_components
= glsl_get_vector_elements(ptr
->type
->type
);
2185 nir_intrinsic_set_write_mask(atomic
, (1 << atomic
->num_components
) - 1);
2186 atomic
->src
[0] = nir_src_for_ssa(vtn_ssa_value(b
, w
[4])->def
);
2187 atomic
->src
[1] = nir_src_for_ssa(index
);
2188 atomic
->src
[2] = nir_src_for_ssa(offset
);
2191 case SpvOpAtomicExchange
:
2192 case SpvOpAtomicCompareExchange
:
2193 case SpvOpAtomicCompareExchangeWeak
:
2194 case SpvOpAtomicIIncrement
:
2195 case SpvOpAtomicIDecrement
:
2196 case SpvOpAtomicIAdd
:
2197 case SpvOpAtomicISub
:
2198 case SpvOpAtomicSMin
:
2199 case SpvOpAtomicUMin
:
2200 case SpvOpAtomicSMax
:
2201 case SpvOpAtomicUMax
:
2202 case SpvOpAtomicAnd
:
2204 case SpvOpAtomicXor
:
2205 atomic
->src
[0] = nir_src_for_ssa(index
);
2206 atomic
->src
[1] = nir_src_for_ssa(offset
);
2207 fill_common_atomic_sources(b
, opcode
, w
, &atomic
->src
[2]);
2211 unreachable("Invalid SPIR-V atomic");
2215 if (opcode
!= SpvOpAtomicStore
) {
2216 struct vtn_type
*type
= vtn_value(b
, w
[1], vtn_value_type_type
)->type
;
2218 nir_ssa_dest_init(&atomic
->instr
, &atomic
->dest
,
2219 glsl_get_vector_elements(type
->type
),
2220 glsl_get_bit_size(type
->type
), NULL
);
2222 struct vtn_value
*val
= vtn_push_value(b
, w
[2], vtn_value_type_ssa
);
2223 val
->ssa
= rzalloc(b
, struct vtn_ssa_value
);
2224 val
->ssa
->def
= &atomic
->dest
.ssa
;
2225 val
->ssa
->type
= type
->type
;
2228 nir_builder_instr_insert(&b
->nb
, &atomic
->instr
);
2231 static nir_alu_instr
*
2232 create_vec(nir_shader
*shader
, unsigned num_components
, unsigned bit_size
)
2235 switch (num_components
) {
2236 case 1: op
= nir_op_fmov
; break;
2237 case 2: op
= nir_op_vec2
; break;
2238 case 3: op
= nir_op_vec3
; break;
2239 case 4: op
= nir_op_vec4
; break;
2240 default: unreachable("bad vector size");
2243 nir_alu_instr
*vec
= nir_alu_instr_create(shader
, op
);
2244 nir_ssa_dest_init(&vec
->instr
, &vec
->dest
.dest
, num_components
,
2246 vec
->dest
.write_mask
= (1 << num_components
) - 1;
2251 struct vtn_ssa_value
*
2252 vtn_ssa_transpose(struct vtn_builder
*b
, struct vtn_ssa_value
*src
)
2254 if (src
->transposed
)
2255 return src
->transposed
;
2257 struct vtn_ssa_value
*dest
=
2258 vtn_create_ssa_value(b
, glsl_transposed_type(src
->type
));
2260 for (unsigned i
= 0; i
< glsl_get_matrix_columns(dest
->type
); i
++) {
2261 nir_alu_instr
*vec
= create_vec(b
->shader
,
2262 glsl_get_matrix_columns(src
->type
),
2263 glsl_get_bit_size(src
->type
));
2264 if (glsl_type_is_vector_or_scalar(src
->type
)) {
2265 vec
->src
[0].src
= nir_src_for_ssa(src
->def
);
2266 vec
->src
[0].swizzle
[0] = i
;
2268 for (unsigned j
= 0; j
< glsl_get_matrix_columns(src
->type
); j
++) {
2269 vec
->src
[j
].src
= nir_src_for_ssa(src
->elems
[j
]->def
);
2270 vec
->src
[j
].swizzle
[0] = i
;
2273 nir_builder_instr_insert(&b
->nb
, &vec
->instr
);
2274 dest
->elems
[i
]->def
= &vec
->dest
.dest
.ssa
;
2277 dest
->transposed
= src
;
2283 vtn_vector_extract(struct vtn_builder
*b
, nir_ssa_def
*src
, unsigned index
)
2285 unsigned swiz
[4] = { index
};
2286 return nir_swizzle(&b
->nb
, src
, swiz
, 1, true);
2290 vtn_vector_insert(struct vtn_builder
*b
, nir_ssa_def
*src
, nir_ssa_def
*insert
,
2293 nir_alu_instr
*vec
= create_vec(b
->shader
, src
->num_components
,
2296 for (unsigned i
= 0; i
< src
->num_components
; i
++) {
2298 vec
->src
[i
].src
= nir_src_for_ssa(insert
);
2300 vec
->src
[i
].src
= nir_src_for_ssa(src
);
2301 vec
->src
[i
].swizzle
[0] = i
;
2305 nir_builder_instr_insert(&b
->nb
, &vec
->instr
);
2307 return &vec
->dest
.dest
.ssa
;
2311 vtn_vector_extract_dynamic(struct vtn_builder
*b
, nir_ssa_def
*src
,
2314 nir_ssa_def
*dest
= vtn_vector_extract(b
, src
, 0);
2315 for (unsigned i
= 1; i
< src
->num_components
; i
++)
2316 dest
= nir_bcsel(&b
->nb
, nir_ieq(&b
->nb
, index
, nir_imm_int(&b
->nb
, i
)),
2317 vtn_vector_extract(b
, src
, i
), dest
);
2323 vtn_vector_insert_dynamic(struct vtn_builder
*b
, nir_ssa_def
*src
,
2324 nir_ssa_def
*insert
, nir_ssa_def
*index
)
2326 nir_ssa_def
*dest
= vtn_vector_insert(b
, src
, insert
, 0);
2327 for (unsigned i
= 1; i
< src
->num_components
; i
++)
2328 dest
= nir_bcsel(&b
->nb
, nir_ieq(&b
->nb
, index
, nir_imm_int(&b
->nb
, i
)),
2329 vtn_vector_insert(b
, src
, insert
, i
), dest
);
2334 static nir_ssa_def
*
2335 vtn_vector_shuffle(struct vtn_builder
*b
, unsigned num_components
,
2336 nir_ssa_def
*src0
, nir_ssa_def
*src1
,
2337 const uint32_t *indices
)
2339 nir_alu_instr
*vec
= create_vec(b
->shader
, num_components
, src0
->bit_size
);
2341 for (unsigned i
= 0; i
< num_components
; i
++) {
2342 uint32_t index
= indices
[i
];
2343 if (index
== 0xffffffff) {
2345 nir_src_for_ssa(nir_ssa_undef(&b
->nb
, 1, src0
->bit_size
));
2346 } else if (index
< src0
->num_components
) {
2347 vec
->src
[i
].src
= nir_src_for_ssa(src0
);
2348 vec
->src
[i
].swizzle
[0] = index
;
2350 vec
->src
[i
].src
= nir_src_for_ssa(src1
);
2351 vec
->src
[i
].swizzle
[0] = index
- src0
->num_components
;
2355 nir_builder_instr_insert(&b
->nb
, &vec
->instr
);
2357 return &vec
->dest
.dest
.ssa
;
2361 * Concatentates a number of vectors/scalars together to produce a vector
2363 static nir_ssa_def
*
2364 vtn_vector_construct(struct vtn_builder
*b
, unsigned num_components
,
2365 unsigned num_srcs
, nir_ssa_def
**srcs
)
2367 nir_alu_instr
*vec
= create_vec(b
->shader
, num_components
,
2370 /* From the SPIR-V 1.1 spec for OpCompositeConstruct:
2372 * "When constructing a vector, there must be at least two Constituent
2375 assert(num_srcs
>= 2);
2377 unsigned dest_idx
= 0;
2378 for (unsigned i
= 0; i
< num_srcs
; i
++) {
2379 nir_ssa_def
*src
= srcs
[i
];
2380 assert(dest_idx
+ src
->num_components
<= num_components
);
2381 for (unsigned j
= 0; j
< src
->num_components
; j
++) {
2382 vec
->src
[dest_idx
].src
= nir_src_for_ssa(src
);
2383 vec
->src
[dest_idx
].swizzle
[0] = j
;
2388 /* From the SPIR-V 1.1 spec for OpCompositeConstruct:
2390 * "When constructing a vector, the total number of components in all
2391 * the operands must equal the number of components in Result Type."
2393 assert(dest_idx
== num_components
);
2395 nir_builder_instr_insert(&b
->nb
, &vec
->instr
);
2397 return &vec
->dest
.dest
.ssa
;
2400 static struct vtn_ssa_value
*
2401 vtn_composite_copy(void *mem_ctx
, struct vtn_ssa_value
*src
)
2403 struct vtn_ssa_value
*dest
= rzalloc(mem_ctx
, struct vtn_ssa_value
);
2404 dest
->type
= src
->type
;
2406 if (glsl_type_is_vector_or_scalar(src
->type
)) {
2407 dest
->def
= src
->def
;
2409 unsigned elems
= glsl_get_length(src
->type
);
2411 dest
->elems
= ralloc_array(mem_ctx
, struct vtn_ssa_value
*, elems
);
2412 for (unsigned i
= 0; i
< elems
; i
++)
2413 dest
->elems
[i
] = vtn_composite_copy(mem_ctx
, src
->elems
[i
]);
2419 static struct vtn_ssa_value
*
2420 vtn_composite_insert(struct vtn_builder
*b
, struct vtn_ssa_value
*src
,
2421 struct vtn_ssa_value
*insert
, const uint32_t *indices
,
2422 unsigned num_indices
)
2424 struct vtn_ssa_value
*dest
= vtn_composite_copy(b
, src
);
2426 struct vtn_ssa_value
*cur
= dest
;
2428 for (i
= 0; i
< num_indices
- 1; i
++) {
2429 cur
= cur
->elems
[indices
[i
]];
2432 if (glsl_type_is_vector_or_scalar(cur
->type
)) {
2433 /* According to the SPIR-V spec, OpCompositeInsert may work down to
2434 * the component granularity. In that case, the last index will be
2435 * the index to insert the scalar into the vector.
2438 cur
->def
= vtn_vector_insert(b
, cur
->def
, insert
->def
, indices
[i
]);
2440 cur
->elems
[indices
[i
]] = insert
;
2446 static struct vtn_ssa_value
*
2447 vtn_composite_extract(struct vtn_builder
*b
, struct vtn_ssa_value
*src
,
2448 const uint32_t *indices
, unsigned num_indices
)
2450 struct vtn_ssa_value
*cur
= src
;
2451 for (unsigned i
= 0; i
< num_indices
; i
++) {
2452 if (glsl_type_is_vector_or_scalar(cur
->type
)) {
2453 assert(i
== num_indices
- 1);
2454 /* According to the SPIR-V spec, OpCompositeExtract may work down to
2455 * the component granularity. The last index will be the index of the
2456 * vector to extract.
2459 struct vtn_ssa_value
*ret
= rzalloc(b
, struct vtn_ssa_value
);
2460 ret
->type
= glsl_scalar_type(glsl_get_base_type(cur
->type
));
2461 ret
->def
= vtn_vector_extract(b
, cur
->def
, indices
[i
]);
2464 cur
= cur
->elems
[indices
[i
]];
2472 vtn_handle_composite(struct vtn_builder
*b
, SpvOp opcode
,
2473 const uint32_t *w
, unsigned count
)
2475 struct vtn_value
*val
= vtn_push_value(b
, w
[2], vtn_value_type_ssa
);
2476 const struct glsl_type
*type
=
2477 vtn_value(b
, w
[1], vtn_value_type_type
)->type
->type
;
2478 val
->ssa
= vtn_create_ssa_value(b
, type
);
2481 case SpvOpVectorExtractDynamic
:
2482 val
->ssa
->def
= vtn_vector_extract_dynamic(b
, vtn_ssa_value(b
, w
[3])->def
,
2483 vtn_ssa_value(b
, w
[4])->def
);
2486 case SpvOpVectorInsertDynamic
:
2487 val
->ssa
->def
= vtn_vector_insert_dynamic(b
, vtn_ssa_value(b
, w
[3])->def
,
2488 vtn_ssa_value(b
, w
[4])->def
,
2489 vtn_ssa_value(b
, w
[5])->def
);
2492 case SpvOpVectorShuffle
:
2493 val
->ssa
->def
= vtn_vector_shuffle(b
, glsl_get_vector_elements(type
),
2494 vtn_ssa_value(b
, w
[3])->def
,
2495 vtn_ssa_value(b
, w
[4])->def
,
2499 case SpvOpCompositeConstruct
: {
2500 unsigned elems
= count
- 3;
2501 if (glsl_type_is_vector_or_scalar(type
)) {
2502 nir_ssa_def
*srcs
[4];
2503 for (unsigned i
= 0; i
< elems
; i
++)
2504 srcs
[i
] = vtn_ssa_value(b
, w
[3 + i
])->def
;
2506 vtn_vector_construct(b
, glsl_get_vector_elements(type
),
2509 val
->ssa
->elems
= ralloc_array(b
, struct vtn_ssa_value
*, elems
);
2510 for (unsigned i
= 0; i
< elems
; i
++)
2511 val
->ssa
->elems
[i
] = vtn_ssa_value(b
, w
[3 + i
]);
2515 case SpvOpCompositeExtract
:
2516 val
->ssa
= vtn_composite_extract(b
, vtn_ssa_value(b
, w
[3]),
2520 case SpvOpCompositeInsert
:
2521 val
->ssa
= vtn_composite_insert(b
, vtn_ssa_value(b
, w
[4]),
2522 vtn_ssa_value(b
, w
[3]),
2526 case SpvOpCopyObject
:
2527 val
->ssa
= vtn_composite_copy(b
, vtn_ssa_value(b
, w
[3]));
2531 unreachable("unknown composite operation");
2536 vtn_handle_barrier(struct vtn_builder
*b
, SpvOp opcode
,
2537 const uint32_t *w
, unsigned count
)
2539 nir_intrinsic_op intrinsic_op
;
2541 case SpvOpEmitVertex
:
2542 case SpvOpEmitStreamVertex
:
2543 intrinsic_op
= nir_intrinsic_emit_vertex
;
2545 case SpvOpEndPrimitive
:
2546 case SpvOpEndStreamPrimitive
:
2547 intrinsic_op
= nir_intrinsic_end_primitive
;
2549 case SpvOpMemoryBarrier
:
2550 intrinsic_op
= nir_intrinsic_memory_barrier
;
2552 case SpvOpControlBarrier
:
2553 intrinsic_op
= nir_intrinsic_barrier
;
2556 unreachable("unknown barrier instruction");
2559 nir_intrinsic_instr
*intrin
=
2560 nir_intrinsic_instr_create(b
->shader
, intrinsic_op
);
2562 if (opcode
== SpvOpEmitStreamVertex
|| opcode
== SpvOpEndStreamPrimitive
)
2563 nir_intrinsic_set_stream_id(intrin
, w
[1]);
2565 nir_builder_instr_insert(&b
->nb
, &intrin
->instr
);
2569 gl_primitive_from_spv_execution_mode(SpvExecutionMode mode
)
2572 case SpvExecutionModeInputPoints
:
2573 case SpvExecutionModeOutputPoints
:
2574 return 0; /* GL_POINTS */
2575 case SpvExecutionModeInputLines
:
2576 return 1; /* GL_LINES */
2577 case SpvExecutionModeInputLinesAdjacency
:
2578 return 0x000A; /* GL_LINE_STRIP_ADJACENCY_ARB */
2579 case SpvExecutionModeTriangles
:
2580 return 4; /* GL_TRIANGLES */
2581 case SpvExecutionModeInputTrianglesAdjacency
:
2582 return 0x000C; /* GL_TRIANGLES_ADJACENCY_ARB */
2583 case SpvExecutionModeQuads
:
2584 return 7; /* GL_QUADS */
2585 case SpvExecutionModeIsolines
:
2586 return 0x8E7A; /* GL_ISOLINES */
2587 case SpvExecutionModeOutputLineStrip
:
2588 return 3; /* GL_LINE_STRIP */
2589 case SpvExecutionModeOutputTriangleStrip
:
2590 return 5; /* GL_TRIANGLE_STRIP */
2592 assert(!"Invalid primitive type");
2598 vertices_in_from_spv_execution_mode(SpvExecutionMode mode
)
2601 case SpvExecutionModeInputPoints
:
2603 case SpvExecutionModeInputLines
:
2605 case SpvExecutionModeInputLinesAdjacency
:
2607 case SpvExecutionModeTriangles
:
2609 case SpvExecutionModeInputTrianglesAdjacency
:
2612 assert(!"Invalid GS input mode");
2617 static gl_shader_stage
2618 stage_for_execution_model(SpvExecutionModel model
)
2621 case SpvExecutionModelVertex
:
2622 return MESA_SHADER_VERTEX
;
2623 case SpvExecutionModelTessellationControl
:
2624 return MESA_SHADER_TESS_CTRL
;
2625 case SpvExecutionModelTessellationEvaluation
:
2626 return MESA_SHADER_TESS_EVAL
;
2627 case SpvExecutionModelGeometry
:
2628 return MESA_SHADER_GEOMETRY
;
2629 case SpvExecutionModelFragment
:
2630 return MESA_SHADER_FRAGMENT
;
2631 case SpvExecutionModelGLCompute
:
2632 return MESA_SHADER_COMPUTE
;
2634 unreachable("Unsupported execution model");
2638 #define spv_check_supported(name, cap) do { \
2639 if (!(b->ext && b->ext->name)) \
2640 vtn_warn("Unsupported SPIR-V capability: %s", \
2641 spirv_capability_to_string(cap)); \
2645 vtn_handle_preamble_instruction(struct vtn_builder
*b
, SpvOp opcode
,
2646 const uint32_t *w
, unsigned count
)
2650 case SpvOpSourceExtension
:
2651 case SpvOpSourceContinued
:
2652 case SpvOpExtension
:
2653 /* Unhandled, but these are for debug so that's ok. */
2656 case SpvOpCapability
: {
2657 SpvCapability cap
= w
[1];
2659 case SpvCapabilityMatrix
:
2660 case SpvCapabilityShader
:
2661 case SpvCapabilityGeometry
:
2662 case SpvCapabilityGeometryPointSize
:
2663 case SpvCapabilityUniformBufferArrayDynamicIndexing
:
2664 case SpvCapabilitySampledImageArrayDynamicIndexing
:
2665 case SpvCapabilityStorageBufferArrayDynamicIndexing
:
2666 case SpvCapabilityStorageImageArrayDynamicIndexing
:
2667 case SpvCapabilityImageRect
:
2668 case SpvCapabilitySampledRect
:
2669 case SpvCapabilitySampled1D
:
2670 case SpvCapabilityImage1D
:
2671 case SpvCapabilitySampledCubeArray
:
2672 case SpvCapabilitySampledBuffer
:
2673 case SpvCapabilityImageBuffer
:
2674 case SpvCapabilityImageQuery
:
2675 case SpvCapabilityDerivativeControl
:
2676 case SpvCapabilityInterpolationFunction
:
2677 case SpvCapabilityMultiViewport
:
2678 case SpvCapabilitySampleRateShading
:
2679 case SpvCapabilityClipDistance
:
2680 case SpvCapabilityCullDistance
:
2681 case SpvCapabilityInputAttachment
:
2682 case SpvCapabilityImageGatherExtended
:
2683 case SpvCapabilityStorageImageExtendedFormats
:
2686 case SpvCapabilityGeometryStreams
:
2687 case SpvCapabilityLinkage
:
2688 case SpvCapabilityVector16
:
2689 case SpvCapabilityFloat16Buffer
:
2690 case SpvCapabilityFloat16
:
2691 case SpvCapabilityInt64Atomics
:
2692 case SpvCapabilityAtomicStorage
:
2693 case SpvCapabilityInt16
:
2694 case SpvCapabilityStorageImageMultisample
:
2695 case SpvCapabilityImageCubeArray
:
2696 case SpvCapabilityInt8
:
2697 case SpvCapabilitySparseResidency
:
2698 case SpvCapabilityMinLod
:
2699 case SpvCapabilityTransformFeedback
:
2700 vtn_warn("Unsupported SPIR-V capability: %s",
2701 spirv_capability_to_string(cap
));
2704 case SpvCapabilityFloat64
:
2705 spv_check_supported(float64
, cap
);
2707 case SpvCapabilityInt64
:
2708 spv_check_supported(int64
, cap
);
2711 case SpvCapabilityAddresses
:
2712 case SpvCapabilityKernel
:
2713 case SpvCapabilityImageBasic
:
2714 case SpvCapabilityImageReadWrite
:
2715 case SpvCapabilityImageMipmap
:
2716 case SpvCapabilityPipes
:
2717 case SpvCapabilityGroups
:
2718 case SpvCapabilityDeviceEnqueue
:
2719 case SpvCapabilityLiteralSampler
:
2720 case SpvCapabilityGenericPointer
:
2721 vtn_warn("Unsupported OpenCL-style SPIR-V capability: %s",
2722 spirv_capability_to_string(cap
));
2725 case SpvCapabilityImageMSArray
:
2726 spv_check_supported(image_ms_array
, cap
);
2729 case SpvCapabilityTessellation
:
2730 case SpvCapabilityTessellationPointSize
:
2731 spv_check_supported(tessellation
, cap
);
2734 case SpvCapabilityDrawParameters
:
2735 spv_check_supported(draw_parameters
, cap
);
2738 case SpvCapabilityStorageImageReadWithoutFormat
:
2739 spv_check_supported(image_read_without_format
, cap
);
2742 case SpvCapabilityStorageImageWriteWithoutFormat
:
2743 spv_check_supported(image_write_without_format
, cap
);
2746 case SpvCapabilityMultiView
:
2747 spv_check_supported(multiview
, cap
);
2751 unreachable("Unhandled capability");
2756 case SpvOpExtInstImport
:
2757 vtn_handle_extension(b
, opcode
, w
, count
);
2760 case SpvOpMemoryModel
:
2761 assert(w
[1] == SpvAddressingModelLogical
);
2762 assert(w
[2] == SpvMemoryModelGLSL450
);
2765 case SpvOpEntryPoint
: {
2766 struct vtn_value
*entry_point
= &b
->values
[w
[2]];
2767 /* Let this be a name label regardless */
2768 unsigned name_words
;
2769 entry_point
->name
= vtn_string_literal(b
, &w
[3], count
- 3, &name_words
);
2771 if (strcmp(entry_point
->name
, b
->entry_point_name
) != 0 ||
2772 stage_for_execution_model(w
[1]) != b
->entry_point_stage
)
2775 assert(b
->entry_point
== NULL
);
2776 b
->entry_point
= entry_point
;
2781 vtn_push_value(b
, w
[1], vtn_value_type_string
)->str
=
2782 vtn_string_literal(b
, &w
[2], count
- 2, NULL
);
2786 b
->values
[w
[1]].name
= vtn_string_literal(b
, &w
[2], count
- 2, NULL
);
2789 case SpvOpMemberName
:
2793 case SpvOpExecutionMode
:
2794 case SpvOpDecorationGroup
:
2796 case SpvOpMemberDecorate
:
2797 case SpvOpGroupDecorate
:
2798 case SpvOpGroupMemberDecorate
:
2799 vtn_handle_decoration(b
, opcode
, w
, count
);
2803 return false; /* End of preamble */
2810 vtn_handle_execution_mode(struct vtn_builder
*b
, struct vtn_value
*entry_point
,
2811 const struct vtn_decoration
*mode
, void *data
)
2813 assert(b
->entry_point
== entry_point
);
2815 switch(mode
->exec_mode
) {
2816 case SpvExecutionModeOriginUpperLeft
:
2817 case SpvExecutionModeOriginLowerLeft
:
2818 b
->origin_upper_left
=
2819 (mode
->exec_mode
== SpvExecutionModeOriginUpperLeft
);
2822 case SpvExecutionModeEarlyFragmentTests
:
2823 assert(b
->shader
->stage
== MESA_SHADER_FRAGMENT
);
2824 b
->shader
->info
.fs
.early_fragment_tests
= true;
2827 case SpvExecutionModeInvocations
:
2828 assert(b
->shader
->stage
== MESA_SHADER_GEOMETRY
);
2829 b
->shader
->info
.gs
.invocations
= MAX2(1, mode
->literals
[0]);
2832 case SpvExecutionModeDepthReplacing
:
2833 assert(b
->shader
->stage
== MESA_SHADER_FRAGMENT
);
2834 b
->shader
->info
.fs
.depth_layout
= FRAG_DEPTH_LAYOUT_ANY
;
2836 case SpvExecutionModeDepthGreater
:
2837 assert(b
->shader
->stage
== MESA_SHADER_FRAGMENT
);
2838 b
->shader
->info
.fs
.depth_layout
= FRAG_DEPTH_LAYOUT_GREATER
;
2840 case SpvExecutionModeDepthLess
:
2841 assert(b
->shader
->stage
== MESA_SHADER_FRAGMENT
);
2842 b
->shader
->info
.fs
.depth_layout
= FRAG_DEPTH_LAYOUT_LESS
;
2844 case SpvExecutionModeDepthUnchanged
:
2845 assert(b
->shader
->stage
== MESA_SHADER_FRAGMENT
);
2846 b
->shader
->info
.fs
.depth_layout
= FRAG_DEPTH_LAYOUT_UNCHANGED
;
2849 case SpvExecutionModeLocalSize
:
2850 assert(b
->shader
->stage
== MESA_SHADER_COMPUTE
);
2851 b
->shader
->info
.cs
.local_size
[0] = mode
->literals
[0];
2852 b
->shader
->info
.cs
.local_size
[1] = mode
->literals
[1];
2853 b
->shader
->info
.cs
.local_size
[2] = mode
->literals
[2];
2855 case SpvExecutionModeLocalSizeHint
:
2856 break; /* Nothing to do with this */
2858 case SpvExecutionModeOutputVertices
:
2859 if (b
->shader
->stage
== MESA_SHADER_TESS_CTRL
||
2860 b
->shader
->stage
== MESA_SHADER_TESS_EVAL
) {
2861 b
->shader
->info
.tess
.tcs_vertices_out
= mode
->literals
[0];
2863 assert(b
->shader
->stage
== MESA_SHADER_GEOMETRY
);
2864 b
->shader
->info
.gs
.vertices_out
= mode
->literals
[0];
2868 case SpvExecutionModeInputPoints
:
2869 case SpvExecutionModeInputLines
:
2870 case SpvExecutionModeInputLinesAdjacency
:
2871 case SpvExecutionModeTriangles
:
2872 case SpvExecutionModeInputTrianglesAdjacency
:
2873 case SpvExecutionModeQuads
:
2874 case SpvExecutionModeIsolines
:
2875 if (b
->shader
->stage
== MESA_SHADER_TESS_CTRL
||
2876 b
->shader
->stage
== MESA_SHADER_TESS_EVAL
) {
2877 b
->shader
->info
.tess
.primitive_mode
=
2878 gl_primitive_from_spv_execution_mode(mode
->exec_mode
);
2880 assert(b
->shader
->stage
== MESA_SHADER_GEOMETRY
);
2881 b
->shader
->info
.gs
.vertices_in
=
2882 vertices_in_from_spv_execution_mode(mode
->exec_mode
);
2886 case SpvExecutionModeOutputPoints
:
2887 case SpvExecutionModeOutputLineStrip
:
2888 case SpvExecutionModeOutputTriangleStrip
:
2889 assert(b
->shader
->stage
== MESA_SHADER_GEOMETRY
);
2890 b
->shader
->info
.gs
.output_primitive
=
2891 gl_primitive_from_spv_execution_mode(mode
->exec_mode
);
2894 case SpvExecutionModeSpacingEqual
:
2895 assert(b
->shader
->stage
== MESA_SHADER_TESS_CTRL
||
2896 b
->shader
->stage
== MESA_SHADER_TESS_EVAL
);
2897 b
->shader
->info
.tess
.spacing
= TESS_SPACING_EQUAL
;
2899 case SpvExecutionModeSpacingFractionalEven
:
2900 assert(b
->shader
->stage
== MESA_SHADER_TESS_CTRL
||
2901 b
->shader
->stage
== MESA_SHADER_TESS_EVAL
);
2902 b
->shader
->info
.tess
.spacing
= TESS_SPACING_FRACTIONAL_EVEN
;
2904 case SpvExecutionModeSpacingFractionalOdd
:
2905 assert(b
->shader
->stage
== MESA_SHADER_TESS_CTRL
||
2906 b
->shader
->stage
== MESA_SHADER_TESS_EVAL
);
2907 b
->shader
->info
.tess
.spacing
= TESS_SPACING_FRACTIONAL_ODD
;
2909 case SpvExecutionModeVertexOrderCw
:
2910 assert(b
->shader
->stage
== MESA_SHADER_TESS_CTRL
||
2911 b
->shader
->stage
== MESA_SHADER_TESS_EVAL
);
2912 /* Vulkan's notion of CCW seems to match the hardware backends,
2913 * but be the opposite of OpenGL. Currently NIR follows GL semantics,
2914 * so we set it backwards here.
2916 b
->shader
->info
.tess
.ccw
= true;
2918 case SpvExecutionModeVertexOrderCcw
:
2919 assert(b
->shader
->stage
== MESA_SHADER_TESS_CTRL
||
2920 b
->shader
->stage
== MESA_SHADER_TESS_EVAL
);
2921 /* Backwards; see above */
2922 b
->shader
->info
.tess
.ccw
= false;
2924 case SpvExecutionModePointMode
:
2925 assert(b
->shader
->stage
== MESA_SHADER_TESS_CTRL
||
2926 b
->shader
->stage
== MESA_SHADER_TESS_EVAL
);
2927 b
->shader
->info
.tess
.point_mode
= true;
2930 case SpvExecutionModePixelCenterInteger
:
2931 b
->pixel_center_integer
= true;
2934 case SpvExecutionModeXfb
:
2935 assert(!"Unhandled execution mode");
2938 case SpvExecutionModeVecTypeHint
:
2939 case SpvExecutionModeContractionOff
:
2943 unreachable("Unhandled execution mode");
2948 vtn_handle_variable_or_type_instruction(struct vtn_builder
*b
, SpvOp opcode
,
2949 const uint32_t *w
, unsigned count
)
2953 case SpvOpSourceContinued
:
2954 case SpvOpSourceExtension
:
2955 case SpvOpExtension
:
2956 case SpvOpCapability
:
2957 case SpvOpExtInstImport
:
2958 case SpvOpMemoryModel
:
2959 case SpvOpEntryPoint
:
2960 case SpvOpExecutionMode
:
2963 case SpvOpMemberName
:
2964 case SpvOpDecorationGroup
:
2966 case SpvOpMemberDecorate
:
2967 case SpvOpGroupDecorate
:
2968 case SpvOpGroupMemberDecorate
:
2969 assert(!"Invalid opcode types and variables section");
2975 case SpvOpTypeFloat
:
2976 case SpvOpTypeVector
:
2977 case SpvOpTypeMatrix
:
2978 case SpvOpTypeImage
:
2979 case SpvOpTypeSampler
:
2980 case SpvOpTypeSampledImage
:
2981 case SpvOpTypeArray
:
2982 case SpvOpTypeRuntimeArray
:
2983 case SpvOpTypeStruct
:
2984 case SpvOpTypeOpaque
:
2985 case SpvOpTypePointer
:
2986 case SpvOpTypeFunction
:
2987 case SpvOpTypeEvent
:
2988 case SpvOpTypeDeviceEvent
:
2989 case SpvOpTypeReserveId
:
2990 case SpvOpTypeQueue
:
2992 vtn_handle_type(b
, opcode
, w
, count
);
2995 case SpvOpConstantTrue
:
2996 case SpvOpConstantFalse
:
2998 case SpvOpConstantComposite
:
2999 case SpvOpConstantSampler
:
3000 case SpvOpConstantNull
:
3001 case SpvOpSpecConstantTrue
:
3002 case SpvOpSpecConstantFalse
:
3003 case SpvOpSpecConstant
:
3004 case SpvOpSpecConstantComposite
:
3005 case SpvOpSpecConstantOp
:
3006 vtn_handle_constant(b
, opcode
, w
, count
);
3011 vtn_handle_variables(b
, opcode
, w
, count
);
3015 return false; /* End of preamble */
3022 vtn_handle_body_instruction(struct vtn_builder
*b
, SpvOp opcode
,
3023 const uint32_t *w
, unsigned count
)
3029 case SpvOpLoopMerge
:
3030 case SpvOpSelectionMerge
:
3031 /* This is handled by cfg pre-pass and walk_blocks */
3035 struct vtn_value
*val
= vtn_push_value(b
, w
[2], vtn_value_type_undef
);
3036 val
->type
= vtn_value(b
, w
[1], vtn_value_type_type
)->type
;
3041 vtn_handle_extension(b
, opcode
, w
, count
);
3047 case SpvOpCopyMemory
:
3048 case SpvOpCopyMemorySized
:
3049 case SpvOpAccessChain
:
3050 case SpvOpInBoundsAccessChain
:
3051 case SpvOpArrayLength
:
3052 vtn_handle_variables(b
, opcode
, w
, count
);
3055 case SpvOpFunctionCall
:
3056 vtn_handle_function_call(b
, opcode
, w
, count
);
3059 case SpvOpSampledImage
:
3061 case SpvOpImageSampleImplicitLod
:
3062 case SpvOpImageSampleExplicitLod
:
3063 case SpvOpImageSampleDrefImplicitLod
:
3064 case SpvOpImageSampleDrefExplicitLod
:
3065 case SpvOpImageSampleProjImplicitLod
:
3066 case SpvOpImageSampleProjExplicitLod
:
3067 case SpvOpImageSampleProjDrefImplicitLod
:
3068 case SpvOpImageSampleProjDrefExplicitLod
:
3069 case SpvOpImageFetch
:
3070 case SpvOpImageGather
:
3071 case SpvOpImageDrefGather
:
3072 case SpvOpImageQuerySizeLod
:
3073 case SpvOpImageQueryLod
:
3074 case SpvOpImageQueryLevels
:
3075 case SpvOpImageQuerySamples
:
3076 vtn_handle_texture(b
, opcode
, w
, count
);
3079 case SpvOpImageRead
:
3080 case SpvOpImageWrite
:
3081 case SpvOpImageTexelPointer
:
3082 vtn_handle_image(b
, opcode
, w
, count
);
3085 case SpvOpImageQuerySize
: {
3086 struct vtn_pointer
*image
=
3087 vtn_value(b
, w
[3], vtn_value_type_pointer
)->pointer
;
3088 if (image
->mode
== vtn_variable_mode_image
) {
3089 vtn_handle_image(b
, opcode
, w
, count
);
3091 assert(image
->mode
== vtn_variable_mode_sampler
);
3092 vtn_handle_texture(b
, opcode
, w
, count
);
3097 case SpvOpAtomicLoad
:
3098 case SpvOpAtomicExchange
:
3099 case SpvOpAtomicCompareExchange
:
3100 case SpvOpAtomicCompareExchangeWeak
:
3101 case SpvOpAtomicIIncrement
:
3102 case SpvOpAtomicIDecrement
:
3103 case SpvOpAtomicIAdd
:
3104 case SpvOpAtomicISub
:
3105 case SpvOpAtomicSMin
:
3106 case SpvOpAtomicUMin
:
3107 case SpvOpAtomicSMax
:
3108 case SpvOpAtomicUMax
:
3109 case SpvOpAtomicAnd
:
3111 case SpvOpAtomicXor
: {
3112 struct vtn_value
*pointer
= vtn_untyped_value(b
, w
[3]);
3113 if (pointer
->value_type
== vtn_value_type_image_pointer
) {
3114 vtn_handle_image(b
, opcode
, w
, count
);
3116 assert(pointer
->value_type
== vtn_value_type_pointer
);
3117 vtn_handle_ssbo_or_shared_atomic(b
, opcode
, w
, count
);
3122 case SpvOpAtomicStore
: {
3123 struct vtn_value
*pointer
= vtn_untyped_value(b
, w
[1]);
3124 if (pointer
->value_type
== vtn_value_type_image_pointer
) {
3125 vtn_handle_image(b
, opcode
, w
, count
);
3127 assert(pointer
->value_type
== vtn_value_type_pointer
);
3128 vtn_handle_ssbo_or_shared_atomic(b
, opcode
, w
, count
);
3138 case SpvOpConvertFToU
:
3139 case SpvOpConvertFToS
:
3140 case SpvOpConvertSToF
:
3141 case SpvOpConvertUToF
:
3145 case SpvOpQuantizeToF16
:
3146 case SpvOpConvertPtrToU
:
3147 case SpvOpConvertUToPtr
:
3148 case SpvOpPtrCastToGeneric
:
3149 case SpvOpGenericCastToPtr
:
3155 case SpvOpSignBitSet
:
3156 case SpvOpLessOrGreater
:
3158 case SpvOpUnordered
:
3173 case SpvOpVectorTimesScalar
:
3175 case SpvOpIAddCarry
:
3176 case SpvOpISubBorrow
:
3177 case SpvOpUMulExtended
:
3178 case SpvOpSMulExtended
:
3179 case SpvOpShiftRightLogical
:
3180 case SpvOpShiftRightArithmetic
:
3181 case SpvOpShiftLeftLogical
:
3182 case SpvOpLogicalEqual
:
3183 case SpvOpLogicalNotEqual
:
3184 case SpvOpLogicalOr
:
3185 case SpvOpLogicalAnd
:
3186 case SpvOpLogicalNot
:
3187 case SpvOpBitwiseOr
:
3188 case SpvOpBitwiseXor
:
3189 case SpvOpBitwiseAnd
:
3192 case SpvOpFOrdEqual
:
3193 case SpvOpFUnordEqual
:
3194 case SpvOpINotEqual
:
3195 case SpvOpFOrdNotEqual
:
3196 case SpvOpFUnordNotEqual
:
3197 case SpvOpULessThan
:
3198 case SpvOpSLessThan
:
3199 case SpvOpFOrdLessThan
:
3200 case SpvOpFUnordLessThan
:
3201 case SpvOpUGreaterThan
:
3202 case SpvOpSGreaterThan
:
3203 case SpvOpFOrdGreaterThan
:
3204 case SpvOpFUnordGreaterThan
:
3205 case SpvOpULessThanEqual
:
3206 case SpvOpSLessThanEqual
:
3207 case SpvOpFOrdLessThanEqual
:
3208 case SpvOpFUnordLessThanEqual
:
3209 case SpvOpUGreaterThanEqual
:
3210 case SpvOpSGreaterThanEqual
:
3211 case SpvOpFOrdGreaterThanEqual
:
3212 case SpvOpFUnordGreaterThanEqual
:
3218 case SpvOpFwidthFine
:
3219 case SpvOpDPdxCoarse
:
3220 case SpvOpDPdyCoarse
:
3221 case SpvOpFwidthCoarse
:
3222 case SpvOpBitFieldInsert
:
3223 case SpvOpBitFieldSExtract
:
3224 case SpvOpBitFieldUExtract
:
3225 case SpvOpBitReverse
:
3227 case SpvOpTranspose
:
3228 case SpvOpOuterProduct
:
3229 case SpvOpMatrixTimesScalar
:
3230 case SpvOpVectorTimesMatrix
:
3231 case SpvOpMatrixTimesVector
:
3232 case SpvOpMatrixTimesMatrix
:
3233 vtn_handle_alu(b
, opcode
, w
, count
);
3236 case SpvOpVectorExtractDynamic
:
3237 case SpvOpVectorInsertDynamic
:
3238 case SpvOpVectorShuffle
:
3239 case SpvOpCompositeConstruct
:
3240 case SpvOpCompositeExtract
:
3241 case SpvOpCompositeInsert
:
3242 case SpvOpCopyObject
:
3243 vtn_handle_composite(b
, opcode
, w
, count
);
3246 case SpvOpEmitVertex
:
3247 case SpvOpEndPrimitive
:
3248 case SpvOpEmitStreamVertex
:
3249 case SpvOpEndStreamPrimitive
:
3250 case SpvOpControlBarrier
:
3251 case SpvOpMemoryBarrier
:
3252 vtn_handle_barrier(b
, opcode
, w
, count
);
3256 unreachable("Unhandled opcode");
3263 spirv_to_nir(const uint32_t *words
, size_t word_count
,
3264 struct nir_spirv_specialization
*spec
, unsigned num_spec
,
3265 gl_shader_stage stage
, const char *entry_point_name
,
3266 const struct nir_spirv_supported_extensions
*ext
,
3267 const nir_shader_compiler_options
*options
)
3269 const uint32_t *word_end
= words
+ word_count
;
3271 /* Handle the SPIR-V header (first 4 dwords) */
3272 assert(word_count
> 5);
3274 assert(words
[0] == SpvMagicNumber
);
3275 assert(words
[1] >= 0x10000);
3276 /* words[2] == generator magic */
3277 unsigned value_id_bound
= words
[3];
3278 assert(words
[4] == 0);
3282 /* Initialize the stn_builder object */
3283 struct vtn_builder
*b
= rzalloc(NULL
, struct vtn_builder
);
3284 b
->value_id_bound
= value_id_bound
;
3285 b
->values
= rzalloc_array(b
, struct vtn_value
, value_id_bound
);
3286 exec_list_make_empty(&b
->functions
);
3287 b
->entry_point_stage
= stage
;
3288 b
->entry_point_name
= entry_point_name
;
3291 /* Handle all the preamble instructions */
3292 words
= vtn_foreach_instruction(b
, words
, word_end
,
3293 vtn_handle_preamble_instruction
);
3295 if (b
->entry_point
== NULL
) {
3296 assert(!"Entry point not found");
3301 b
->shader
= nir_shader_create(NULL
, stage
, options
, NULL
);
3303 /* Set shader info defaults */
3304 b
->shader
->info
.gs
.invocations
= 1;
3306 /* Parse execution modes */
3307 vtn_foreach_execution_mode(b
, b
->entry_point
,
3308 vtn_handle_execution_mode
, NULL
);
3310 b
->specializations
= spec
;
3311 b
->num_specializations
= num_spec
;
3313 /* Handle all variable, type, and constant instructions */
3314 words
= vtn_foreach_instruction(b
, words
, word_end
,
3315 vtn_handle_variable_or_type_instruction
);
3317 vtn_build_cfg(b
, words
, word_end
);
3319 foreach_list_typed(struct vtn_function
, func
, node
, &b
->functions
) {
3320 b
->impl
= func
->impl
;
3321 b
->const_table
= _mesa_hash_table_create(b
, _mesa_hash_pointer
,
3322 _mesa_key_pointer_equal
);
3324 vtn_function_emit(b
, func
, vtn_handle_body_instruction
);
3327 assert(b
->entry_point
->value_type
== vtn_value_type_function
);
3328 nir_function
*entry_point
= b
->entry_point
->func
->impl
->function
;
3329 assert(entry_point
);