projects / mesa.git / blob
? search:
summary | shortlog | log | commit | commitdiff | tree
history | raw | HEAD
glsl,nir: Switch the enum representing shader image formats to PIPE_FORMAT.
[mesa.git] / src / compiler / spirv / spirv_to_nir.c
1 /*
2 * Copyright © 2015 Intel Corporation
3 *
4 * Permission is hereby granted, free of charge, to any person obtaining a
5 * copy of this software and associated documentation files (the "Software"),
6 * to deal in the Software without restriction, including without limitation
7 * the rights to use, copy, modify, merge, publish, distribute, sublicense,
8 * and/or sell copies of the Software, and to permit persons to whom the
9 * Software is furnished to do so, subject to the following conditions:
10 *
11 * The above copyright notice and this permission notice (including the next
12 * paragraph) shall be included in all copies or substantial portions of the
13 * Software.
14 *
15 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
16 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
17 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
18 * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
19 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
20 * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
21 * IN THE SOFTWARE.
22 *
23 * Authors:
24 * Jason Ekstrand (jason@jlekstrand.net)
25 *
26 */
27
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 "nir/nir_deref.h"
33 #include "spirv_info.h"
34
35 #include "util/format/u_format.h"
36 #include "util/u_math.h"
37
38 #include <stdio.h>
39
40 void
41 vtn_log(struct vtn_builder *b, enum nir_spirv_debug_level level,
42 size_t spirv_offset, const char *message)
43 {
44 if (b->options->debug.func) {
45 b->options->debug.func(b->options->debug.private_data,
46 level, spirv_offset, message);
47 }
48
49 #ifndef NDEBUG
50 if (level >= NIR_SPIRV_DEBUG_LEVEL_WARNING)
51 fprintf(stderr, "%s\n", message);
52 #endif
53 }
54
55 void
56 vtn_logf(struct vtn_builder *b, enum nir_spirv_debug_level level,
57 size_t spirv_offset, const char *fmt, ...)
58 {
59 va_list args;
60 char *msg;
61
62 va_start(args, fmt);
63 msg = ralloc_vasprintf(NULL, fmt, args);
64 va_end(args);
65
66 vtn_log(b, level, spirv_offset, msg);
67
68 ralloc_free(msg);
69 }
70
71 static void
72 vtn_log_err(struct vtn_builder *b,
73 enum nir_spirv_debug_level level, const char *prefix,
74 const char *file, unsigned line,
75 const char *fmt, va_list args)
76 {
77 char *msg;
78
79 msg = ralloc_strdup(NULL, prefix);
80
81 #ifndef NDEBUG
82 ralloc_asprintf_append(&msg, " In file %s:%u\n", file, line);
83 #endif
84
85 ralloc_asprintf_append(&msg, " ");
86
87 ralloc_vasprintf_append(&msg, fmt, args);
88
89 ralloc_asprintf_append(&msg, "\n %zu bytes into the SPIR-V binary",
90 b->spirv_offset);
91
92 if (b->file) {
93 ralloc_asprintf_append(&msg,
94 "\n in SPIR-V source file %s, line %d, col %d",
95 b->file, b->line, b->col);
96 }
97
98 vtn_log(b, level, b->spirv_offset, msg);
99
100 ralloc_free(msg);
101 }
102
103 static void
104 vtn_dump_shader(struct vtn_builder *b, const char *path, const char *prefix)
105 {
106 static int idx = 0;
107
108 char filename[1024];
109 int len = snprintf(filename, sizeof(filename), "%s/%s-%d.spirv",
110 path, prefix, idx++);
111 if (len < 0 || len >= sizeof(filename))
112 return;
113
114 FILE *f = fopen(filename, "w");
115 if (f == NULL)
116 return;
117
118 fwrite(b->spirv, sizeof(*b->spirv), b->spirv_word_count, f);
119 fclose(f);
120
121 vtn_info("SPIR-V shader dumped to %s", filename);
122 }
123
124 void
125 _vtn_warn(struct vtn_builder *b, const char *file, unsigned line,
126 const char *fmt, ...)
127 {
128 va_list args;
129
130 va_start(args, fmt);
131 vtn_log_err(b, NIR_SPIRV_DEBUG_LEVEL_WARNING, "SPIR-V WARNING:\n",
132 file, line, fmt, args);
133 va_end(args);
134 }
135
136 void
137 _vtn_err(struct vtn_builder *b, const char *file, unsigned line,
138 const char *fmt, ...)
139 {
140 va_list args;
141
142 va_start(args, fmt);
143 vtn_log_err(b, NIR_SPIRV_DEBUG_LEVEL_ERROR, "SPIR-V ERROR:\n",
144 file, line, fmt, args);
145 va_end(args);
146 }
147
148 void
149 _vtn_fail(struct vtn_builder *b, const char *file, unsigned line,
150 const char *fmt, ...)
151 {
152 va_list args;
153
154 va_start(args, fmt);
155 vtn_log_err(b, NIR_SPIRV_DEBUG_LEVEL_ERROR, "SPIR-V parsing FAILED:\n",
156 file, line, fmt, args);
157 va_end(args);
158
159 const char *dump_path = getenv("MESA_SPIRV_FAIL_DUMP_PATH");
160 if (dump_path)
161 vtn_dump_shader(b, dump_path, "fail");
162
163 longjmp(b->fail_jump, 1);
164 }
165
166 struct spec_constant_value {
167 bool is_double;
168 union {
169 uint32_t data32;
170 uint64_t data64;
171 };
172 };
173
174 static struct vtn_ssa_value *
175 vtn_undef_ssa_value(struct vtn_builder *b, const struct glsl_type *type)
176 {
177 struct vtn_ssa_value *val = rzalloc(b, struct vtn_ssa_value);
178 val->type = type;
179
180 if (glsl_type_is_vector_or_scalar(type)) {
181 unsigned num_components = glsl_get_vector_elements(val->type);
182 unsigned bit_size = glsl_get_bit_size(val->type);
183 val->def = nir_ssa_undef(&b->nb, num_components, bit_size);
184 } else {
185 unsigned elems = glsl_get_length(val->type);
186 val->elems = ralloc_array(b, struct vtn_ssa_value *, elems);
187 if (glsl_type_is_matrix(type)) {
188 const struct glsl_type *elem_type =
189 glsl_vector_type(glsl_get_base_type(type),
190 glsl_get_vector_elements(type));
191
192 for (unsigned i = 0; i < elems; i++)
193 val->elems[i] = vtn_undef_ssa_value(b, elem_type);
194 } else if (glsl_type_is_array(type)) {
195 const struct glsl_type *elem_type = glsl_get_array_element(type);
196 for (unsigned i = 0; i < elems; i++)
197 val->elems[i] = vtn_undef_ssa_value(b, elem_type);
198 } else {
199 for (unsigned i = 0; i < elems; i++) {
200 const struct glsl_type *elem_type = glsl_get_struct_field(type, i);
201 val->elems[i] = vtn_undef_ssa_value(b, elem_type);
202 }
203 }
204 }
205
206 return val;
207 }
208
209 static struct vtn_ssa_value *
210 vtn_const_ssa_value(struct vtn_builder *b, nir_constant *constant,
211 const struct glsl_type *type)
212 {
213 struct hash_entry *entry = _mesa_hash_table_search(b->const_table, constant);
214
215 if (entry)
216 return entry->data;
217
218 struct vtn_ssa_value *val = rzalloc(b, struct vtn_ssa_value);
219 val->type = type;
220
221 switch (glsl_get_base_type(type)) {
222 case GLSL_TYPE_INT:
223 case GLSL_TYPE_UINT:
224 case GLSL_TYPE_INT16:
225 case GLSL_TYPE_UINT16:
226 case GLSL_TYPE_UINT8:
227 case GLSL_TYPE_INT8:
228 case GLSL_TYPE_INT64:
229 case GLSL_TYPE_UINT64:
230 case GLSL_TYPE_BOOL:
231 case GLSL_TYPE_FLOAT:
232 case GLSL_TYPE_FLOAT16:
233 case GLSL_TYPE_DOUBLE: {
234 int bit_size = glsl_get_bit_size(type);
235 if (glsl_type_is_vector_or_scalar(type)) {
236 unsigned num_components = glsl_get_vector_elements(val->type);
237 nir_load_const_instr *load =
238 nir_load_const_instr_create(b->shader, num_components, bit_size);
239
240 memcpy(load->value, constant->values,
241 sizeof(nir_const_value) * load->def.num_components);
242
243 nir_instr_insert_before_cf_list(&b->nb.impl->body, &load->instr);
244 val->def = &load->def;
245 } else {
246 assert(glsl_type_is_matrix(type));
247 unsigned columns = glsl_get_matrix_columns(val->type);
248 val->elems = ralloc_array(b, struct vtn_ssa_value *, columns);
249 const struct glsl_type *column_type = glsl_get_column_type(val->type);
250 for (unsigned i = 0; i < columns; i++)
251 val->elems[i] = vtn_const_ssa_value(b, constant->elements[i],
252 column_type);
253 }
254 break;
255 }
256
257 case GLSL_TYPE_ARRAY: {
258 unsigned elems = glsl_get_length(val->type);
259 val->elems = ralloc_array(b, struct vtn_ssa_value *, elems);
260 const struct glsl_type *elem_type = glsl_get_array_element(val->type);
261 for (unsigned i = 0; i < elems; i++)
262 val->elems[i] = vtn_const_ssa_value(b, constant->elements[i],
263 elem_type);
264 break;
265 }
266
267 case GLSL_TYPE_STRUCT: {
268 unsigned elems = glsl_get_length(val->type);
269 val->elems = ralloc_array(b, struct vtn_ssa_value *, elems);
270 for (unsigned i = 0; i < elems; i++) {
271 const struct glsl_type *elem_type =
272 glsl_get_struct_field(val->type, i);
273 val->elems[i] = vtn_const_ssa_value(b, constant->elements[i],
274 elem_type);
275 }
276 break;
277 }
278
279 default:
280 vtn_fail("bad constant type");
281 }
282
283 return val;
284 }
285
286 struct vtn_ssa_value *
287 vtn_ssa_value(struct vtn_builder *b, uint32_t value_id)
288 {
289 struct vtn_value *val = vtn_untyped_value(b, value_id);
290 switch (val->value_type) {
291 case vtn_value_type_undef:
292 return vtn_undef_ssa_value(b, val->type->type);
293
294 case vtn_value_type_constant:
295 return vtn_const_ssa_value(b, val->constant, val->type->type);
296
297 case vtn_value_type_ssa:
298 return val->ssa;
299
300 case vtn_value_type_pointer:
301 vtn_assert(val->pointer->ptr_type && val->pointer->ptr_type->type);
302 struct vtn_ssa_value *ssa =
303 vtn_create_ssa_value(b, val->pointer->ptr_type->type);
304 ssa->def = vtn_pointer_to_ssa(b, val->pointer);
305 return ssa;
306
307 default:
308 vtn_fail("Invalid type for an SSA value");
309 }
310 }
311
312 static char *
313 vtn_string_literal(struct vtn_builder *b, const uint32_t *words,
314 unsigned word_count, unsigned *words_used)
315 {
316 char *dup = ralloc_strndup(b, (char *)words, word_count * sizeof(*words));
317 if (words_used) {
318 /* Ammount of space taken by the string (including the null) */
319 unsigned len = strlen(dup) + 1;
320 *words_used = DIV_ROUND_UP(len, sizeof(*words));
321 }
322 return dup;
323 }
324
325 const uint32_t *
326 vtn_foreach_instruction(struct vtn_builder *b, const uint32_t *start,
327 const uint32_t *end, vtn_instruction_handler handler)
328 {
329 b->file = NULL;
330 b->line = -1;
331 b->col = -1;
332
333 const uint32_t *w = start;
334 while (w < end) {
335 SpvOp opcode = w[0] & SpvOpCodeMask;
336 unsigned count = w[0] >> SpvWordCountShift;
337 vtn_assert(count >= 1 && w + count <= end);
338
339 b->spirv_offset = (uint8_t *)w - (uint8_t *)b->spirv;
340
341 switch (opcode) {
342 case SpvOpNop:
343 break; /* Do nothing */
344
345 case SpvOpLine:
346 b->file = vtn_value(b, w[1], vtn_value_type_string)->str;
347 b->line = w[2];
348 b->col = w[3];
349 break;
350
351 case SpvOpNoLine:
352 b->file = NULL;
353 b->line = -1;
354 b->col = -1;
355 break;
356
357 default:
358 if (!handler(b, opcode, w, count))
359 return w;
360 break;
361 }
362
363 w += count;
364 }
365
366 b->spirv_offset = 0;
367 b->file = NULL;
368 b->line = -1;
369 b->col = -1;
370
371 assert(w == end);
372 return w;
373 }
374
375 static bool
376 vtn_handle_non_semantic_instruction(struct vtn_builder *b, SpvOp ext_opcode,
377 const uint32_t *w, unsigned count)
378 {
379 /* Do nothing. */
380 return true;
381 }
382
383 static void
384 vtn_handle_extension(struct vtn_builder *b, SpvOp opcode,
385 const uint32_t *w, unsigned count)
386 {
387 const char *ext = (const char *)&w[2];
388 switch (opcode) {
389 case SpvOpExtInstImport: {
390 struct vtn_value *val = vtn_push_value(b, w[1], vtn_value_type_extension);
391 if (strcmp(ext, "GLSL.std.450") == 0) {
392 val->ext_handler = vtn_handle_glsl450_instruction;
393 } else if ((strcmp(ext, "SPV_AMD_gcn_shader") == 0)
394 && (b->options && b->options->caps.amd_gcn_shader)) {
395 val->ext_handler = vtn_handle_amd_gcn_shader_instruction;
396 } else if ((strcmp(ext, "SPV_AMD_shader_ballot") == 0)
397 && (b->options && b->options->caps.amd_shader_ballot)) {
398 val->ext_handler = vtn_handle_amd_shader_ballot_instruction;
399 } else if ((strcmp(ext, "SPV_AMD_shader_trinary_minmax") == 0)
400 && (b->options && b->options->caps.amd_trinary_minmax)) {
401 val->ext_handler = vtn_handle_amd_shader_trinary_minmax_instruction;
402 } else if ((strcmp(ext, "SPV_AMD_shader_explicit_vertex_parameter") == 0)
403 && (b->options && b->options->caps.amd_shader_explicit_vertex_parameter)) {
404 val->ext_handler = vtn_handle_amd_shader_explicit_vertex_parameter_instruction;
405 } else if (strcmp(ext, "OpenCL.std") == 0) {
406 val->ext_handler = vtn_handle_opencl_instruction;
407 } else if (strstr(ext, "NonSemantic.") == ext) {
408 val->ext_handler = vtn_handle_non_semantic_instruction;
409 } else {
410 vtn_fail("Unsupported extension: %s", ext);
411 }
412 break;
413 }
414
415 case SpvOpExtInst: {
416 struct vtn_value *val = vtn_value(b, w[3], vtn_value_type_extension);
417 bool handled = val->ext_handler(b, w[4], w, count);
418 vtn_assert(handled);
419 break;
420 }
421
422 default:
423 vtn_fail_with_opcode("Unhandled opcode", opcode);
424 }
425 }
426
427 static void
428 _foreach_decoration_helper(struct vtn_builder *b,
429 struct vtn_value *base_value,
430 int parent_member,
431 struct vtn_value *value,
432 vtn_decoration_foreach_cb cb, void *data)
433 {
434 for (struct vtn_decoration *dec = value->decoration; dec; dec = dec->next) {
435 int member;
436 if (dec->scope == VTN_DEC_DECORATION) {
437 member = parent_member;
438 } else if (dec->scope >= VTN_DEC_STRUCT_MEMBER0) {
439 vtn_fail_if(value->value_type != vtn_value_type_type ||
440 value->type->base_type != vtn_base_type_struct,
441 "OpMemberDecorate and OpGroupMemberDecorate are only "
442 "allowed on OpTypeStruct");
443 /* This means we haven't recursed yet */
444 assert(value == base_value);
445
446 member = dec->scope - VTN_DEC_STRUCT_MEMBER0;
447
448 vtn_fail_if(member >= base_value->type->length,
449 "OpMemberDecorate specifies member %d but the "
450 "OpTypeStruct has only %u members",
451 member, base_value->type->length);
452 } else {
453 /* Not a decoration */
454 assert(dec->scope == VTN_DEC_EXECUTION_MODE);
455 continue;
456 }
457
458 if (dec->group) {
459 assert(dec->group->value_type == vtn_value_type_decoration_group);
460 _foreach_decoration_helper(b, base_value, member, dec->group,
461 cb, data);
462 } else {
463 cb(b, base_value, member, dec, data);
464 }
465 }
466 }
467
468 /** Iterates (recursively if needed) over all of the decorations on a value
469 *
470 * This function iterates over all of the decorations applied to a given
471 * value. If it encounters a decoration group, it recurses into the group
472 * and iterates over all of those decorations as well.
473 */
474 void
475 vtn_foreach_decoration(struct vtn_builder *b, struct vtn_value *value,
476 vtn_decoration_foreach_cb cb, void *data)
477 {
478 _foreach_decoration_helper(b, value, -1, value, cb, data);
479 }
480
481 void
482 vtn_foreach_execution_mode(struct vtn_builder *b, struct vtn_value *value,
483 vtn_execution_mode_foreach_cb cb, void *data)
484 {
485 for (struct vtn_decoration *dec = value->decoration; dec; dec = dec->next) {
486 if (dec->scope != VTN_DEC_EXECUTION_MODE)
487 continue;
488
489 assert(dec->group == NULL);
490 cb(b, value, dec, data);
491 }
492 }
493
494 void
495 vtn_handle_decoration(struct vtn_builder *b, SpvOp opcode,
496 const uint32_t *w, unsigned count)
497 {
498 const uint32_t *w_end = w + count;
499 const uint32_t target = w[1];
500 w += 2;
501
502 switch (opcode) {
503 case SpvOpDecorationGroup:
504 vtn_push_value(b, target, vtn_value_type_decoration_group);
505 break;
506
507 case SpvOpDecorate:
508 case SpvOpDecorateId:
509 case SpvOpMemberDecorate:
510 case SpvOpDecorateString:
511 case SpvOpMemberDecorateString:
512 case SpvOpExecutionMode:
513 case SpvOpExecutionModeId: {
514 struct vtn_value *val = vtn_untyped_value(b, target);
515
516 struct vtn_decoration *dec = rzalloc(b, struct vtn_decoration);
517 switch (opcode) {
518 case SpvOpDecorate:
519 case SpvOpDecorateId:
520 case SpvOpDecorateString:
521 dec->scope = VTN_DEC_DECORATION;
522 break;
523 case SpvOpMemberDecorate:
524 case SpvOpMemberDecorateString:
525 dec->scope = VTN_DEC_STRUCT_MEMBER0 + *(w++);
526 vtn_fail_if(dec->scope < VTN_DEC_STRUCT_MEMBER0, /* overflow */
527 "Member argument of OpMemberDecorate too large");
528 break;
529 case SpvOpExecutionMode:
530 case SpvOpExecutionModeId:
531 dec->scope = VTN_DEC_EXECUTION_MODE;
532 break;
533 default:
534 unreachable("Invalid decoration opcode");
535 }
536 dec->decoration = *(w++);
537 dec->operands = w;
538
539 /* Link into the list */
540 dec->next = val->decoration;
541 val->decoration = dec;
542 break;
543 }
544
545 case SpvOpGroupMemberDecorate:
546 case SpvOpGroupDecorate: {
547 struct vtn_value *group =
548 vtn_value(b, target, vtn_value_type_decoration_group);
549
550 for (; w < w_end; w++) {
551 struct vtn_value *val = vtn_untyped_value(b, *w);
552 struct vtn_decoration *dec = rzalloc(b, struct vtn_decoration);
553
554 dec->group = group;
555 if (opcode == SpvOpGroupDecorate) {
556 dec->scope = VTN_DEC_DECORATION;
557 } else {
558 dec->scope = VTN_DEC_STRUCT_MEMBER0 + *(++w);
559 vtn_fail_if(dec->scope < 0, /* Check for overflow */
560 "Member argument of OpGroupMemberDecorate too large");
561 }
562
563 /* Link into the list */
564 dec->next = val->decoration;
565 val->decoration = dec;
566 }
567 break;
568 }
569
570 default:
571 unreachable("Unhandled opcode");
572 }
573 }
574
575 struct member_decoration_ctx {
576 unsigned num_fields;
577 struct glsl_struct_field *fields;
578 struct vtn_type *type;
579 };
580
581 /**
582 * Returns true if the given type contains a struct decorated Block or
583 * BufferBlock
584 */
585 bool
586 vtn_type_contains_block(struct vtn_builder *b, struct vtn_type *type)
587 {
588 switch (type->base_type) {
589 case vtn_base_type_array:
590 return vtn_type_contains_block(b, type->array_element);
591 case vtn_base_type_struct:
592 if (type->block || type->buffer_block)
593 return true;
594 for (unsigned i = 0; i < type->length; i++) {
595 if (vtn_type_contains_block(b, type->members[i]))
596 return true;
597 }
598 return false;
599 default:
600 return false;
601 }
602 }
603
604 /** Returns true if two types are "compatible", i.e. you can do an OpLoad,
605 * OpStore, or OpCopyMemory between them without breaking anything.
606 * Technically, the SPIR-V rules require the exact same type ID but this lets
607 * us internally be a bit looser.
608 */
609 bool
610 vtn_types_compatible(struct vtn_builder *b,
611 struct vtn_type *t1, struct vtn_type *t2)
612 {
613 if (t1->id == t2->id)
614 return true;
615
616 if (t1->base_type != t2->base_type)
617 return false;
618
619 switch (t1->base_type) {
620 case vtn_base_type_void:
621 case vtn_base_type_scalar:
622 case vtn_base_type_vector:
623 case vtn_base_type_matrix:
624 case vtn_base_type_image:
625 case vtn_base_type_sampler:
626 case vtn_base_type_sampled_image:
627 return t1->type == t2->type;
628
629 case vtn_base_type_array:
630 return t1->length == t2->length &&
631 vtn_types_compatible(b, t1->array_element, t2->array_element);
632
633 case vtn_base_type_pointer:
634 return vtn_types_compatible(b, t1->deref, t2->deref);
635
636 case vtn_base_type_struct:
637 if (t1->length != t2->length)
638 return false;
639
640 for (unsigned i = 0; i < t1->length; i++) {
641 if (!vtn_types_compatible(b, t1->members[i], t2->members[i]))
642 return false;
643 }
644 return true;
645
646 case vtn_base_type_function:
647 /* This case shouldn't get hit since you can't copy around function
648 * types. Just require them to be identical.
649 */
650 return false;
651 }
652
653 vtn_fail("Invalid base type");
654 }
655
656 struct vtn_type *
657 vtn_type_without_array(struct vtn_type *type)
658 {
659 while (type->base_type == vtn_base_type_array)
660 type = type->array_element;
661 return type;
662 }
663
664 /* does a shallow copy of a vtn_type */
665
666 static struct vtn_type *
667 vtn_type_copy(struct vtn_builder *b, struct vtn_type *src)
668 {
669 struct vtn_type *dest = ralloc(b, struct vtn_type);
670 *dest = *src;
671
672 switch (src->base_type) {
673 case vtn_base_type_void:
674 case vtn_base_type_scalar:
675 case vtn_base_type_vector:
676 case vtn_base_type_matrix:
677 case vtn_base_type_array:
678 case vtn_base_type_pointer:
679 case vtn_base_type_image:
680 case vtn_base_type_sampler:
681 case vtn_base_type_sampled_image:
682 /* Nothing more to do */
683 break;
684
685 case vtn_base_type_struct:
686 dest->members = ralloc_array(b, struct vtn_type *, src->length);
687 memcpy(dest->members, src->members,
688 src->length * sizeof(src->members[0]));
689
690 dest->offsets = ralloc_array(b, unsigned, src->length);
691 memcpy(dest->offsets, src->offsets,
692 src->length * sizeof(src->offsets[0]));
693 break;
694
695 case vtn_base_type_function:
696 dest->params = ralloc_array(b, struct vtn_type *, src->length);
697 memcpy(dest->params, src->params, src->length * sizeof(src->params[0]));
698 break;
699 }
700
701 return dest;
702 }
703
704 static struct vtn_type *
705 mutable_matrix_member(struct vtn_builder *b, struct vtn_type *type, int member)
706 {
707 type->members[member] = vtn_type_copy(b, type->members[member]);
708 type = type->members[member];
709
710 /* We may have an array of matrices.... Oh, joy! */
711 while (glsl_type_is_array(type->type)) {
712 type->array_element = vtn_type_copy(b, type->array_element);
713 type = type->array_element;
714 }
715
716 vtn_assert(glsl_type_is_matrix(type->type));
717
718 return type;
719 }
720
721 static void
722 vtn_handle_access_qualifier(struct vtn_builder *b, struct vtn_type *type,
723 int member, enum gl_access_qualifier access)
724 {
725 type->members[member] = vtn_type_copy(b, type->members[member]);
726 type = type->members[member];
727
728 type->access |= access;
729 }
730
731 static void
732 array_stride_decoration_cb(struct vtn_builder *b,
733 struct vtn_value *val, int member,
734 const struct vtn_decoration *dec, void *void_ctx)
735 {
736 struct vtn_type *type = val->type;
737
738 if (dec->decoration == SpvDecorationArrayStride) {
739 if (vtn_type_contains_block(b, type)) {
740 vtn_warn("The ArrayStride decoration cannot be applied to an array "
741 "type which contains a structure type decorated Block "
742 "or BufferBlock");
743 /* Ignore the decoration */
744 } else {
745 vtn_fail_if(dec->operands[0] == 0, "ArrayStride must be non-zero");
746 type->stride = dec->operands[0];
747 }
748 }
749 }
750
751 static void
752 struct_member_decoration_cb(struct vtn_builder *b,
753 UNUSED struct vtn_value *val, int member,
754 const struct vtn_decoration *dec, void *void_ctx)
755 {
756 struct member_decoration_ctx *ctx = void_ctx;
757
758 if (member < 0)
759 return;
760
761 assert(member < ctx->num_fields);
762
763 switch (dec->decoration) {
764 case SpvDecorationRelaxedPrecision:
765 case SpvDecorationUniform:
766 case SpvDecorationUniformId:
767 break; /* FIXME: Do nothing with this for now. */
768 case SpvDecorationNonWritable:
769 vtn_handle_access_qualifier(b, ctx->type, member, ACCESS_NON_WRITEABLE);
770 break;
771 case SpvDecorationNonReadable:
772 vtn_handle_access_qualifier(b, ctx->type, member, ACCESS_NON_READABLE);
773 break;
774 case SpvDecorationVolatile:
775 vtn_handle_access_qualifier(b, ctx->type, member, ACCESS_VOLATILE);
776 break;
777 case SpvDecorationCoherent:
778 vtn_handle_access_qualifier(b, ctx->type, member, ACCESS_COHERENT);
779 break;
780 case SpvDecorationNoPerspective:
781 ctx->fields[member].interpolation = INTERP_MODE_NOPERSPECTIVE;
782 break;
783 case SpvDecorationFlat:
784 ctx->fields[member].interpolation = INTERP_MODE_FLAT;
785 break;
786 case SpvDecorationExplicitInterpAMD:
787 ctx->fields[member].interpolation = INTERP_MODE_EXPLICIT;
788 break;
789 case SpvDecorationCentroid:
790 ctx->fields[member].centroid = true;
791 break;
792 case SpvDecorationSample:
793 ctx->fields[member].sample = true;
794 break;
795 case SpvDecorationStream:
796 /* Vulkan only allows one GS stream */
797 vtn_assert(dec->operands[0] == 0);
798 break;
799 case SpvDecorationLocation:
800 ctx->fields[member].location = dec->operands[0];
801 break;
802 case SpvDecorationComponent:
803 break; /* FIXME: What should we do with these? */
804 case SpvDecorationBuiltIn:
805 ctx->type->members[member] = vtn_type_copy(b, ctx->type->members[member]);
806 ctx->type->members[member]->is_builtin = true;
807 ctx->type->members[member]->builtin = dec->operands[0];
808 ctx->type->builtin_block = true;
809 break;
810 case SpvDecorationOffset:
811 ctx->type->offsets[member] = dec->operands[0];
812 ctx->fields[member].offset = dec->operands[0];
813 break;
814 case SpvDecorationMatrixStride:
815 /* Handled as a second pass */
816 break;
817 case SpvDecorationColMajor:
818 break; /* Nothing to do here. Column-major is the default. */
819 case SpvDecorationRowMajor:
820 mutable_matrix_member(b, ctx->type, member)->row_major = true;
821 break;
822
823 case SpvDecorationPatch:
824 break;
825
826 case SpvDecorationSpecId:
827 case SpvDecorationBlock:
828 case SpvDecorationBufferBlock:
829 case SpvDecorationArrayStride:
830 case SpvDecorationGLSLShared:
831 case SpvDecorationGLSLPacked:
832 case SpvDecorationInvariant:
833 case SpvDecorationRestrict:
834 case SpvDecorationAliased:
835 case SpvDecorationConstant:
836 case SpvDecorationIndex:
837 case SpvDecorationBinding:
838 case SpvDecorationDescriptorSet:
839 case SpvDecorationLinkageAttributes:
840 case SpvDecorationNoContraction:
841 case SpvDecorationInputAttachmentIndex:
842 vtn_warn("Decoration not allowed on struct members: %s",
843 spirv_decoration_to_string(dec->decoration));
844 break;
845
846 case SpvDecorationXfbBuffer:
847 case SpvDecorationXfbStride:
848 vtn_warn("Vulkan does not have transform feedback");
849 break;
850
851 case SpvDecorationCPacked:
852 if (b->shader->info.stage != MESA_SHADER_KERNEL)
853 vtn_warn("Decoration only allowed for CL-style kernels: %s",
854 spirv_decoration_to_string(dec->decoration));
855 else
856 ctx->type->packed = true;
857 break;
858
859 case SpvDecorationSaturatedConversion:
860 case SpvDecorationFuncParamAttr:
861 case SpvDecorationFPRoundingMode:
862 case SpvDecorationFPFastMathMode:
863 case SpvDecorationAlignment:
864 if (b->shader->info.stage != MESA_SHADER_KERNEL) {
865 vtn_warn("Decoration only allowed for CL-style kernels: %s",
866 spirv_decoration_to_string(dec->decoration));
867 }
868 break;
869
870 case SpvDecorationUserSemantic:
871 /* User semantic decorations can safely be ignored by the driver. */
872 break;
873
874 default:
875 vtn_fail_with_decoration("Unhandled decoration", dec->decoration);
876 }
877 }
878
879 /** Chases the array type all the way down to the tail and rewrites the
880 * glsl_types to be based off the tail's glsl_type.
881 */
882 static void
883 vtn_array_type_rewrite_glsl_type(struct vtn_type *type)
884 {
885 if (type->base_type != vtn_base_type_array)
886 return;
887
888 vtn_array_type_rewrite_glsl_type(type->array_element);
889
890 type->type = glsl_array_type(type->array_element->type,
891 type->length, type->stride);
892 }
893
894 /* Matrix strides are handled as a separate pass because we need to know
895 * whether the matrix is row-major or not first.
896 */
897 static void
898 struct_member_matrix_stride_cb(struct vtn_builder *b,
899 UNUSED struct vtn_value *val, int member,
900 const struct vtn_decoration *dec,
901 void *void_ctx)
902 {
903 if (dec->decoration != SpvDecorationMatrixStride)
904 return;
905
906 vtn_fail_if(member < 0,
907 "The MatrixStride decoration is only allowed on members "
908 "of OpTypeStruct");
909 vtn_fail_if(dec->operands[0] == 0, "MatrixStride must be non-zero");
910
911 struct member_decoration_ctx *ctx = void_ctx;
912
913 struct vtn_type *mat_type = mutable_matrix_member(b, ctx->type, member);
914 if (mat_type->row_major) {
915 mat_type->array_element = vtn_type_copy(b, mat_type->array_element);
916 mat_type->stride = mat_type->array_element->stride;
917 mat_type->array_element->stride = dec->operands[0];
918
919 mat_type->type = glsl_explicit_matrix_type(mat_type->type,
920 dec->operands[0], true);
921 mat_type->array_element->type = glsl_get_column_type(mat_type->type);
922 } else {
923 vtn_assert(mat_type->array_element->stride > 0);
924 mat_type->stride = dec->operands[0];
925
926 mat_type->type = glsl_explicit_matrix_type(mat_type->type,
927 dec->operands[0], false);
928 }
929
930 /* Now that we've replaced the glsl_type with a properly strided matrix
931 * type, rewrite the member type so that it's an array of the proper kind
932 * of glsl_type.
933 */
934 vtn_array_type_rewrite_glsl_type(ctx->type->members[member]);
935 ctx->fields[member].type = ctx->type->members[member]->type;
936 }
937
938 static void
939 struct_block_decoration_cb(struct vtn_builder *b,
940 struct vtn_value *val, int member,
941 const struct vtn_decoration *dec, void *ctx)
942 {
943 if (member != -1)
944 return;
945
946 struct vtn_type *type = val->type;
947 if (dec->decoration == SpvDecorationBlock)
948 type->block = true;
949 else if (dec->decoration == SpvDecorationBufferBlock)
950 type->buffer_block = true;
951 }
952
953 static void
954 type_decoration_cb(struct vtn_builder *b,
955 struct vtn_value *val, int member,
956 const struct vtn_decoration *dec, UNUSED void *ctx)
957 {
958 struct vtn_type *type = val->type;
959
960 if (member != -1) {
961 /* This should have been handled by OpTypeStruct */
962 assert(val->type->base_type == vtn_base_type_struct);
963 assert(member >= 0 && member < val->type->length);
964 return;
965 }
966
967 switch (dec->decoration) {
968 case SpvDecorationArrayStride:
969 vtn_assert(type->base_type == vtn_base_type_array ||
970 type->base_type == vtn_base_type_pointer);
971 break;
972 case SpvDecorationBlock:
973 vtn_assert(type->base_type == vtn_base_type_struct);
974 vtn_assert(type->block);
975 break;
976 case SpvDecorationBufferBlock:
977 vtn_assert(type->base_type == vtn_base_type_struct);
978 vtn_assert(type->buffer_block);
979 break;
980 case SpvDecorationGLSLShared:
981 case SpvDecorationGLSLPacked:
982 /* Ignore these, since we get explicit offsets anyways */
983 break;
984
985 case SpvDecorationRowMajor:
986 case SpvDecorationColMajor:
987 case SpvDecorationMatrixStride:
988 case SpvDecorationBuiltIn:
989 case SpvDecorationNoPerspective:
990 case SpvDecorationFlat:
991 case SpvDecorationPatch:
992 case SpvDecorationCentroid:
993 case SpvDecorationSample:
994 case SpvDecorationExplicitInterpAMD:
995 case SpvDecorationVolatile:
996 case SpvDecorationCoherent:
997 case SpvDecorationNonWritable:
998 case SpvDecorationNonReadable:
999 case SpvDecorationUniform:
1000 case SpvDecorationUniformId:
1001 case SpvDecorationLocation:
1002 case SpvDecorationComponent:
1003 case SpvDecorationOffset:
1004 case SpvDecorationXfbBuffer:
1005 case SpvDecorationXfbStride:
1006 case SpvDecorationUserSemantic:
1007 vtn_warn("Decoration only allowed for struct members: %s",
1008 spirv_decoration_to_string(dec->decoration));
1009 break;
1010
1011 case SpvDecorationStream:
1012 /* We don't need to do anything here, as stream is filled up when
1013 * aplying the decoration to a variable, just check that if it is not a
1014 * struct member, it should be a struct.
1015 */
1016 vtn_assert(type->base_type == vtn_base_type_struct);
1017 break;
1018
1019 case SpvDecorationRelaxedPrecision:
1020 case SpvDecorationSpecId:
1021 case SpvDecorationInvariant:
1022 case SpvDecorationRestrict:
1023 case SpvDecorationAliased:
1024 case SpvDecorationConstant:
1025 case SpvDecorationIndex:
1026 case SpvDecorationBinding:
1027 case SpvDecorationDescriptorSet:
1028 case SpvDecorationLinkageAttributes:
1029 case SpvDecorationNoContraction:
1030 case SpvDecorationInputAttachmentIndex:
1031 vtn_warn("Decoration not allowed on types: %s",
1032 spirv_decoration_to_string(dec->decoration));
1033 break;
1034
1035 case SpvDecorationCPacked:
1036 if (b->shader->info.stage != MESA_SHADER_KERNEL)
1037 vtn_warn("Decoration only allowed for CL-style kernels: %s",
1038 spirv_decoration_to_string(dec->decoration));
1039 else
1040 type->packed = true;
1041 break;
1042
1043 case SpvDecorationSaturatedConversion:
1044 case SpvDecorationFuncParamAttr:
1045 case SpvDecorationFPRoundingMode:
1046 case SpvDecorationFPFastMathMode:
1047 case SpvDecorationAlignment:
1048 vtn_warn("Decoration only allowed for CL-style kernels: %s",
1049 spirv_decoration_to_string(dec->decoration));
1050 break;
1051
1052 default:
1053 vtn_fail_with_decoration("Unhandled decoration", dec->decoration);
1054 }
1055 }
1056
1057 static unsigned
1058 translate_image_format(struct vtn_builder *b, SpvImageFormat format)
1059 {
1060 switch (format) {
1061 case SpvImageFormatUnknown: return PIPE_FORMAT_NONE;
1062 case SpvImageFormatRgba32f: return PIPE_FORMAT_R32G32B32A32_FLOAT;
1063 case SpvImageFormatRgba16f: return PIPE_FORMAT_R16G16B16A16_FLOAT;
1064 case SpvImageFormatR32f: return PIPE_FORMAT_R32_FLOAT;
1065 case SpvImageFormatRgba8: return PIPE_FORMAT_R8G8B8A8_UNORM;
1066 case SpvImageFormatRgba8Snorm: return PIPE_FORMAT_R8G8B8A8_SNORM;
1067 case SpvImageFormatRg32f: return PIPE_FORMAT_R32G32_FLOAT;
1068 case SpvImageFormatRg16f: return PIPE_FORMAT_R16G16_FLOAT;
1069 case SpvImageFormatR11fG11fB10f: return PIPE_FORMAT_R11G11B10_FLOAT;
1070 case SpvImageFormatR16f: return PIPE_FORMAT_R16_FLOAT;
1071 case SpvImageFormatRgba16: return PIPE_FORMAT_R16G16B16A16_UNORM;
1072 case SpvImageFormatRgb10A2: return PIPE_FORMAT_R10G10B10A2_UNORM;
1073 case SpvImageFormatRg16: return PIPE_FORMAT_R16G16_UNORM;
1074 case SpvImageFormatRg8: return PIPE_FORMAT_R8G8_UNORM;
1075 case SpvImageFormatR16: return PIPE_FORMAT_R16_UNORM;
1076 case SpvImageFormatR8: return PIPE_FORMAT_R8_UNORM;
1077 case SpvImageFormatRgba16Snorm: return PIPE_FORMAT_R16G16B16A16_SNORM;
1078 case SpvImageFormatRg16Snorm: return PIPE_FORMAT_R16G16_SNORM;
1079 case SpvImageFormatRg8Snorm: return PIPE_FORMAT_R8G8_SNORM;
1080 case SpvImageFormatR16Snorm: return PIPE_FORMAT_R16_SNORM;
1081 case SpvImageFormatR8Snorm: return PIPE_FORMAT_R8_SNORM;
1082 case SpvImageFormatRgba32i: return PIPE_FORMAT_R32G32B32A32_SINT;
1083 case SpvImageFormatRgba16i: return PIPE_FORMAT_R16G16B16A16_SINT;
1084 case SpvImageFormatRgba8i: return PIPE_FORMAT_R8G8B8A8_SINT;
1085 case SpvImageFormatR32i: return PIPE_FORMAT_R32_SINT;
1086 case SpvImageFormatRg32i: return PIPE_FORMAT_R32G32_SINT;
1087 case SpvImageFormatRg16i: return PIPE_FORMAT_R16G16_SINT;
1088 case SpvImageFormatRg8i: return PIPE_FORMAT_R8G8_SINT;
1089 case SpvImageFormatR16i: return PIPE_FORMAT_R16_SINT;
1090 case SpvImageFormatR8i: return PIPE_FORMAT_R8_SINT;
1091 case SpvImageFormatRgba32ui: return PIPE_FORMAT_R32G32B32A32_UINT;
1092 case SpvImageFormatRgba16ui: return PIPE_FORMAT_R16G16B16A16_UINT;
1093 case SpvImageFormatRgba8ui: return PIPE_FORMAT_R8G8B8A8_UINT;
1094 case SpvImageFormatR32ui: return PIPE_FORMAT_R32_UINT;
1095 case SpvImageFormatRgb10a2ui: return PIPE_FORMAT_R10G10B10A2_UINT;
1096 case SpvImageFormatRg32ui: return PIPE_FORMAT_R32G32_UINT;
1097 case SpvImageFormatRg16ui: return PIPE_FORMAT_R16G16_UINT;
1098 case SpvImageFormatRg8ui: return PIPE_FORMAT_R8G8_UINT;
1099 case SpvImageFormatR16ui: return PIPE_FORMAT_R16_UINT;
1100 case SpvImageFormatR8ui: return PIPE_FORMAT_R8_UINT;
1101 default:
1102 vtn_fail("Invalid image format: %s (%u)",
1103 spirv_imageformat_to_string(format), format);
1104 }
1105 }
1106
1107 static void
1108 vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
1109 const uint32_t *w, unsigned count)
1110 {
1111 struct vtn_value *val = NULL;
1112
1113 /* In order to properly handle forward declarations, we have to defer
1114 * allocation for pointer types.
1115 */
1116 if (opcode != SpvOpTypePointer && opcode != SpvOpTypeForwardPointer) {
1117 val = vtn_push_value(b, w[1], vtn_value_type_type);
1118 vtn_fail_if(val->type != NULL,
1119 "Only pointers can have forward declarations");
1120 val->type = rzalloc(b, struct vtn_type);
1121 val->type->id = w[1];
1122 }
1123
1124 switch (opcode) {
1125 case SpvOpTypeVoid:
1126 val->type->base_type = vtn_base_type_void;
1127 val->type->type = glsl_void_type();
1128 break;
1129 case SpvOpTypeBool:
1130 val->type->base_type = vtn_base_type_scalar;
1131 val->type->type = glsl_bool_type();
1132 val->type->length = 1;
1133 break;
1134 case SpvOpTypeInt: {
1135 int bit_size = w[2];
1136 const bool signedness = w[3];
1137 val->type->base_type = vtn_base_type_scalar;
1138 switch (bit_size) {
1139 case 64:
1140 val->type->type = (signedness ? glsl_int64_t_type() : glsl_uint64_t_type());
1141 break;
1142 case 32:
1143 val->type->type = (signedness ? glsl_int_type() : glsl_uint_type());
1144 break;
1145 case 16:
1146 val->type->type = (signedness ? glsl_int16_t_type() : glsl_uint16_t_type());
1147 break;
1148 case 8:
1149 val->type->type = (signedness ? glsl_int8_t_type() : glsl_uint8_t_type());
1150 break;
1151 default:
1152 vtn_fail("Invalid int bit size: %u", bit_size);
1153 }
1154 val->type->length = 1;
1155 break;
1156 }
1157
1158 case SpvOpTypeFloat: {
1159 int bit_size = w[2];
1160 val->type->base_type = vtn_base_type_scalar;
1161 switch (bit_size) {
1162 case 16:
1163 val->type->type = glsl_float16_t_type();
1164 break;
1165 case 32:
1166 val->type->type = glsl_float_type();
1167 break;
1168 case 64:
1169 val->type->type = glsl_double_type();
1170 break;
1171 default:
1172 vtn_fail("Invalid float bit size: %u", bit_size);
1173 }
1174 val->type->length = 1;
1175 break;
1176 }
1177
1178 case SpvOpTypeVector: {
1179 struct vtn_type *base = vtn_value(b, w[2], vtn_value_type_type)->type;
1180 unsigned elems = w[3];
1181
1182 vtn_fail_if(base->base_type != vtn_base_type_scalar,
1183 "Base type for OpTypeVector must be a scalar");
1184 vtn_fail_if((elems < 2 || elems > 4) && (elems != 8) && (elems != 16),
1185 "Invalid component count for OpTypeVector");
1186
1187 val->type->base_type = vtn_base_type_vector;
1188 val->type->type = glsl_vector_type(glsl_get_base_type(base->type), elems);
1189 val->type->length = elems;
1190 val->type->stride = glsl_type_is_boolean(val->type->type)
1191 ? 4 : glsl_get_bit_size(base->type) / 8;
1192 val->type->array_element = base;
1193 break;
1194 }
1195
1196 case SpvOpTypeMatrix: {
1197 struct vtn_type *base = vtn_value(b, w[2], vtn_value_type_type)->type;
1198 unsigned columns = w[3];
1199
1200 vtn_fail_if(base->base_type != vtn_base_type_vector,
1201 "Base type for OpTypeMatrix must be a vector");
1202 vtn_fail_if(columns < 2 || columns > 4,
1203 "Invalid column count for OpTypeMatrix");
1204
1205 val->type->base_type = vtn_base_type_matrix;
1206 val->type->type = glsl_matrix_type(glsl_get_base_type(base->type),
1207 glsl_get_vector_elements(base->type),
1208 columns);
1209 vtn_fail_if(glsl_type_is_error(val->type->type),
1210 "Unsupported base type for OpTypeMatrix");
1211 assert(!glsl_type_is_error(val->type->type));
1212 val->type->length = columns;
1213 val->type->array_element = base;
1214 val->type->row_major = false;
1215 val->type->stride = 0;
1216 break;
1217 }
1218
1219 case SpvOpTypeRuntimeArray:
1220 case SpvOpTypeArray: {
1221 struct vtn_type *array_element =
1222 vtn_value(b, w[2], vtn_value_type_type)->type;
1223
1224 if (opcode == SpvOpTypeRuntimeArray) {
1225 /* A length of 0 is used to denote unsized arrays */
1226 val->type->length = 0;
1227 } else {
1228 val->type->length = vtn_constant_uint(b, w[3]);
1229 }
1230
1231 val->type->base_type = vtn_base_type_array;
1232 val->type->array_element = array_element;
1233 if (b->shader->info.stage == MESA_SHADER_KERNEL)
1234 val->type->stride = glsl_get_cl_size(array_element->type);
1235
1236 vtn_foreach_decoration(b, val, array_stride_decoration_cb, NULL);
1237 val->type->type = glsl_array_type(array_element->type, val->type->length,
1238 val->type->stride);
1239 break;
1240 }
1241
1242 case SpvOpTypeStruct: {
1243 unsigned num_fields = count - 2;
1244 val->type->base_type = vtn_base_type_struct;
1245 val->type->length = num_fields;
1246 val->type->members = ralloc_array(b, struct vtn_type *, num_fields);
1247 val->type->offsets = ralloc_array(b, unsigned, num_fields);
1248 val->type->packed = false;
1249
1250 NIR_VLA(struct glsl_struct_field, fields, count);
1251 for (unsigned i = 0; i < num_fields; i++) {
1252 val->type->members[i] =
1253 vtn_value(b, w[i + 2], vtn_value_type_type)->type;
1254 fields[i] = (struct glsl_struct_field) {
1255 .type = val->type->members[i]->type,
1256 .name = ralloc_asprintf(b, "field%d", i),
1257 .location = -1,
1258 .offset = -1,
1259 };
1260 }
1261
1262 if (b->shader->info.stage == MESA_SHADER_KERNEL) {
1263 unsigned offset = 0;
1264 for (unsigned i = 0; i < num_fields; i++) {
1265 offset = align(offset, glsl_get_cl_alignment(fields[i].type));
1266 fields[i].offset = offset;
1267 offset += glsl_get_cl_size(fields[i].type);
1268 }
1269 }
1270
1271 struct member_decoration_ctx ctx = {
1272 .num_fields = num_fields,
1273 .fields = fields,
1274 .type = val->type
1275 };
1276
1277 vtn_foreach_decoration(b, val, struct_member_decoration_cb, &ctx);
1278 vtn_foreach_decoration(b, val, struct_member_matrix_stride_cb, &ctx);
1279
1280 vtn_foreach_decoration(b, val, struct_block_decoration_cb, NULL);
1281
1282 const char *name = val->name;
1283
1284 if (val->type->block || val->type->buffer_block) {
1285 /* Packing will be ignored since types coming from SPIR-V are
1286 * explicitly laid out.
1287 */
1288 val->type->type = glsl_interface_type(fields, num_fields,
1289 /* packing */ 0, false,
1290 name ? name : "block");
1291 } else {
1292 val->type->type = glsl_struct_type(fields, num_fields,
1293 name ? name : "struct", false);
1294 }
1295 break;
1296 }
1297
1298 case SpvOpTypeFunction: {
1299 val->type->base_type = vtn_base_type_function;
1300 val->type->type = NULL;
1301
1302 val->type->return_type = vtn_value(b, w[2], vtn_value_type_type)->type;
1303
1304 const unsigned num_params = count - 3;
1305 val->type->length = num_params;
1306 val->type->params = ralloc_array(b, struct vtn_type *, num_params);
1307 for (unsigned i = 0; i < count - 3; i++) {
1308 val->type->params[i] =
1309 vtn_value(b, w[i + 3], vtn_value_type_type)->type;
1310 }
1311 break;
1312 }
1313
1314 case SpvOpTypePointer:
1315 case SpvOpTypeForwardPointer: {
1316 /* We can't blindly push the value because it might be a forward
1317 * declaration.
1318 */
1319 val = vtn_untyped_value(b, w[1]);
1320
1321 SpvStorageClass storage_class = w[2];
1322
1323 if (val->value_type == vtn_value_type_invalid) {
1324 val->value_type = vtn_value_type_type;
1325 val->type = rzalloc(b, struct vtn_type);
1326 val->type->id = w[1];
1327 val->type->base_type = vtn_base_type_pointer;
1328 val->type->storage_class = storage_class;
1329
1330 /* These can actually be stored to nir_variables and used as SSA
1331 * values so they need a real glsl_type.
1332 */
1333 enum vtn_variable_mode mode = vtn_storage_class_to_mode(
1334 b, storage_class, NULL, NULL);
1335 val->type->type = nir_address_format_to_glsl_type(
1336 vtn_mode_to_address_format(b, mode));
1337 } else {
1338 vtn_fail_if(val->type->storage_class != storage_class,
1339 "The storage classes of an OpTypePointer and any "
1340 "OpTypeForwardPointers that provide forward "
1341 "declarations of it must match.");
1342 }
1343
1344 if (opcode == SpvOpTypePointer) {
1345 vtn_fail_if(val->type->deref != NULL,
1346 "While OpTypeForwardPointer can be used to provide a "
1347 "forward declaration of a pointer, OpTypePointer can "
1348 "only be used once for a given id.");
1349
1350 val->type->deref = vtn_value(b, w[3], vtn_value_type_type)->type;
1351
1352 /* Only certain storage classes use ArrayStride. The others (in
1353 * particular Workgroup) are expected to be laid out by the driver.
1354 */
1355 switch (storage_class) {
1356 case SpvStorageClassUniform:
1357 case SpvStorageClassPushConstant:
1358 case SpvStorageClassStorageBuffer:
1359 case SpvStorageClassPhysicalStorageBuffer:
1360 vtn_foreach_decoration(b, val, array_stride_decoration_cb, NULL);
1361 break;
1362 default:
1363 /* Nothing to do. */
1364 break;
1365 }
1366
1367 if (b->physical_ptrs) {
1368 switch (storage_class) {
1369 case SpvStorageClassFunction:
1370 case SpvStorageClassWorkgroup:
1371 case SpvStorageClassCrossWorkgroup:
1372 case SpvStorageClassUniformConstant:
1373 val->type->stride = align(glsl_get_cl_size(val->type->deref->type),
1374 glsl_get_cl_alignment(val->type->deref->type));
1375 break;
1376 default:
1377 break;
1378 }
1379 }
1380 }
1381 break;
1382 }
1383
1384 case SpvOpTypeImage: {
1385 val->type->base_type = vtn_base_type_image;
1386
1387 const struct vtn_type *sampled_type =
1388 vtn_value(b, w[2], vtn_value_type_type)->type;
1389
1390 vtn_fail_if(sampled_type->base_type != vtn_base_type_scalar ||
1391 glsl_get_bit_size(sampled_type->type) != 32,
1392 "Sampled type of OpTypeImage must be a 32-bit scalar");
1393
1394 enum glsl_sampler_dim dim;
1395 switch ((SpvDim)w[3]) {
1396 case SpvDim1D: dim = GLSL_SAMPLER_DIM_1D; break;
1397 case SpvDim2D: dim = GLSL_SAMPLER_DIM_2D; break;
1398 case SpvDim3D: dim = GLSL_SAMPLER_DIM_3D; break;
1399 case SpvDimCube: dim = GLSL_SAMPLER_DIM_CUBE; break;
1400 case SpvDimRect: dim = GLSL_SAMPLER_DIM_RECT; break;
1401 case SpvDimBuffer: dim = GLSL_SAMPLER_DIM_BUF; break;
1402 case SpvDimSubpassData: dim = GLSL_SAMPLER_DIM_SUBPASS; break;
1403 default:
1404 vtn_fail("Invalid SPIR-V image dimensionality: %s (%u)",
1405 spirv_dim_to_string((SpvDim)w[3]), w[3]);
1406 }
1407
1408 /* w[4]: as per Vulkan spec "Validation Rules within a Module",
1409 * The “Depth” operand of OpTypeImage is ignored.
1410 */
1411 bool is_array = w[5];
1412 bool multisampled = w[6];
1413 unsigned sampled = w[7];
1414 SpvImageFormat format = w[8];
1415
1416 if (count > 9)
1417 val->type->access_qualifier = w[9];
1418 else
1419 val->type->access_qualifier = SpvAccessQualifierReadWrite;
1420
1421 if (multisampled) {
1422 if (dim == GLSL_SAMPLER_DIM_2D)
1423 dim = GLSL_SAMPLER_DIM_MS;
1424 else if (dim == GLSL_SAMPLER_DIM_SUBPASS)
1425 dim = GLSL_SAMPLER_DIM_SUBPASS_MS;
1426 else
1427 vtn_fail("Unsupported multisampled image type");
1428 }
1429
1430 val->type->image_format = translate_image_format(b, format);
1431
1432 enum glsl_base_type sampled_base_type =
1433 glsl_get_base_type(sampled_type->type);
1434 if (sampled == 1) {
1435 val->type->sampled = true;
1436 val->type->type = glsl_sampler_type(dim, false, is_array,
1437 sampled_base_type);
1438 } else if (sampled == 2) {
1439 val->type->sampled = false;
1440 val->type->type = glsl_image_type(dim, is_array, sampled_base_type);
1441 } else {
1442 vtn_fail("We need to know if the image will be sampled");
1443 }
1444 break;
1445 }
1446
1447 case SpvOpTypeSampledImage:
1448 val->type->base_type = vtn_base_type_sampled_image;
1449 val->type->image = vtn_value(b, w[2], vtn_value_type_type)->type;
1450 val->type->type = val->type->image->type;
1451 break;
1452
1453 case SpvOpTypeSampler:
1454 /* The actual sampler type here doesn't really matter. It gets
1455 * thrown away the moment you combine it with an image. What really
1456 * matters is that it's a sampler type as opposed to an integer type
1457 * so the backend knows what to do.
1458 */
1459 val->type->base_type = vtn_base_type_sampler;
1460 val->type->type = glsl_bare_sampler_type();
1461 break;
1462
1463 case SpvOpTypeOpaque:
1464 case SpvOpTypeEvent:
1465 case SpvOpTypeDeviceEvent:
1466 case SpvOpTypeReserveId:
1467 case SpvOpTypeQueue:
1468 case SpvOpTypePipe:
1469 default:
1470 vtn_fail_with_opcode("Unhandled opcode", opcode);
1471 }
1472
1473 vtn_foreach_decoration(b, val, type_decoration_cb, NULL);
1474
1475 if (val->type->base_type == vtn_base_type_struct &&
1476 (val->type->block || val->type->buffer_block)) {
1477 for (unsigned i = 0; i < val->type->length; i++) {
1478 vtn_fail_if(vtn_type_contains_block(b, val->type->members[i]),
1479 "Block and BufferBlock decorations cannot decorate a "
1480 "structure type that is nested at any level inside "
1481 "another structure type decorated with Block or "
1482 "BufferBlock.");
1483 }
1484 }
1485 }
1486
1487 static nir_constant *
1488 vtn_null_constant(struct vtn_builder *b, struct vtn_type *type)
1489 {
1490 nir_constant *c = rzalloc(b, nir_constant);
1491
1492 switch (type->base_type) {
1493 case vtn_base_type_scalar:
1494 case vtn_base_type_vector:
1495 /* Nothing to do here. It's already initialized to zero */
1496 break;
1497
1498 case vtn_base_type_pointer: {
1499 enum vtn_variable_mode mode = vtn_storage_class_to_mode(
1500 b, type->storage_class, type->deref, NULL);
1501 nir_address_format addr_format = vtn_mode_to_address_format(b, mode);
1502
1503 const nir_const_value *null_value = nir_address_format_null_value(addr_format);
1504 memcpy(c->values, null_value,
1505 sizeof(nir_const_value) * nir_address_format_num_components(addr_format));
1506 break;
1507 }
1508
1509 case vtn_base_type_void:
1510 case vtn_base_type_image:
1511 case vtn_base_type_sampler:
1512 case vtn_base_type_sampled_image:
1513 case vtn_base_type_function:
1514 /* For those we have to return something but it doesn't matter what. */
1515 break;
1516
1517 case vtn_base_type_matrix:
1518 case vtn_base_type_array:
1519 vtn_assert(type->length > 0);
1520 c->num_elements = type->length;
1521 c->elements = ralloc_array(b, nir_constant *, c->num_elements);
1522
1523 c->elements[0] = vtn_null_constant(b, type->array_element);
1524 for (unsigned i = 1; i < c->num_elements; i++)
1525 c->elements[i] = c->elements[0];
1526 break;
1527
1528 case vtn_base_type_struct:
1529 c->num_elements = type->length;
1530 c->elements = ralloc_array(b, nir_constant *, c->num_elements);
1531 for (unsigned i = 0; i < c->num_elements; i++)
1532 c->elements[i] = vtn_null_constant(b, type->members[i]);
1533 break;
1534
1535 default:
1536 vtn_fail("Invalid type for null constant");
1537 }
1538
1539 return c;
1540 }
1541
1542 static void
1543 spec_constant_decoration_cb(struct vtn_builder *b, UNUSED struct vtn_value *val,
1544 ASSERTED int member,
1545 const struct vtn_decoration *dec, void *data)
1546 {
1547 vtn_assert(member == -1);
1548 if (dec->decoration != SpvDecorationSpecId)
1549 return;
1550
1551 struct spec_constant_value *const_value = data;
1552
1553 for (unsigned i = 0; i < b->num_specializations; i++) {
1554 if (b->specializations[i].id == dec->operands[0]) {
1555 if (const_value->is_double)
1556 const_value->data64 = b->specializations[i].data64;
1557 else
1558 const_value->data32 = b->specializations[i].data32;
1559 return;
1560 }
1561 }
1562 }
1563
1564 static uint32_t
1565 get_specialization(struct vtn_builder *b, struct vtn_value *val,
1566 uint32_t const_value)
1567 {
1568 struct spec_constant_value data;
1569 data.is_double = false;
1570 data.data32 = const_value;
1571 vtn_foreach_decoration(b, val, spec_constant_decoration_cb, &data);
1572 return data.data32;
1573 }
1574
1575 static uint64_t
1576 get_specialization64(struct vtn_builder *b, struct vtn_value *val,
1577 uint64_t const_value)
1578 {
1579 struct spec_constant_value data;
1580 data.is_double = true;
1581 data.data64 = const_value;
1582 vtn_foreach_decoration(b, val, spec_constant_decoration_cb, &data);
1583 return data.data64;
1584 }
1585
1586 static void
1587 handle_workgroup_size_decoration_cb(struct vtn_builder *b,
1588 struct vtn_value *val,
1589 ASSERTED int member,
1590 const struct vtn_decoration *dec,
1591 UNUSED void *data)
1592 {
1593 vtn_assert(member == -1);
1594 if (dec->decoration != SpvDecorationBuiltIn ||
1595 dec->operands[0] != SpvBuiltInWorkgroupSize)
1596 return;
1597
1598 vtn_assert(val->type->type == glsl_vector_type(GLSL_TYPE_UINT, 3));
1599 b->workgroup_size_builtin = val;
1600 }
1601
1602 static void
1603 vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
1604 const uint32_t *w, unsigned count)
1605 {
1606 struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_constant);
1607 val->constant = rzalloc(b, nir_constant);
1608 switch (opcode) {
1609 case SpvOpConstantTrue:
1610 case SpvOpConstantFalse:
1611 case SpvOpSpecConstantTrue:
1612 case SpvOpSpecConstantFalse: {
1613 vtn_fail_if(val->type->type != glsl_bool_type(),
1614 "Result type of %s must be OpTypeBool",
1615 spirv_op_to_string(opcode));
1616
1617 uint32_t int_val = (opcode == SpvOpConstantTrue ||
1618 opcode == SpvOpSpecConstantTrue);
1619
1620 if (opcode == SpvOpSpecConstantTrue ||
1621 opcode == SpvOpSpecConstantFalse)
1622 int_val = get_specialization(b, val, int_val);
1623
1624 val->constant->values[0].b = int_val != 0;
1625 break;
1626 }
1627
1628 case SpvOpConstant: {
1629 vtn_fail_if(val->type->base_type != vtn_base_type_scalar,
1630 "Result type of %s must be a scalar",
1631 spirv_op_to_string(opcode));
1632 int bit_size = glsl_get_bit_size(val->type->type);
1633 switch (bit_size) {
1634 case 64:
1635 val->constant->values[0].u64 = vtn_u64_literal(&w[3]);
1636 break;
1637 case 32:
1638 val->constant->values[0].u32 = w[3];
1639 break;
1640 case 16:
1641 val->constant->values[0].u16 = w[3];
1642 break;
1643 case 8:
1644 val->constant->values[0].u8 = w[3];
1645 break;
1646 default:
1647 vtn_fail("Unsupported SpvOpConstant bit size: %u", bit_size);
1648 }
1649 break;
1650 }
1651
1652 case SpvOpSpecConstant: {
1653 vtn_fail_if(val->type->base_type != vtn_base_type_scalar,
1654 "Result type of %s must be a scalar",
1655 spirv_op_to_string(opcode));
1656 int bit_size = glsl_get_bit_size(val->type->type);
1657 switch (bit_size) {
1658 case 64:
1659 val->constant->values[0].u64 =
1660 get_specialization64(b, val, vtn_u64_literal(&w[3]));
1661 break;
1662 case 32:
1663 val->constant->values[0].u32 = get_specialization(b, val, w[3]);
1664 break;
1665 case 16:
1666 val->constant->values[0].u16 = get_specialization(b, val, w[3]);
1667 break;
1668 case 8:
1669 val->constant->values[0].u8 = get_specialization(b, val, w[3]);
1670 break;
1671 default:
1672 vtn_fail("Unsupported SpvOpSpecConstant bit size");
1673 }
1674 break;
1675 }
1676
1677 case SpvOpSpecConstantComposite:
1678 case SpvOpConstantComposite: {
1679 unsigned elem_count = count - 3;
1680 vtn_fail_if(elem_count != val->type->length,
1681 "%s has %u constituents, expected %u",
1682 spirv_op_to_string(opcode), elem_count, val->type->length);
1683
1684 nir_constant **elems = ralloc_array(b, nir_constant *, elem_count);
1685 for (unsigned i = 0; i < elem_count; i++) {
1686 struct vtn_value *val = vtn_untyped_value(b, w[i + 3]);
1687
1688 if (val->value_type == vtn_value_type_constant) {
1689 elems[i] = val->constant;
1690 } else {
1691 vtn_fail_if(val->value_type != vtn_value_type_undef,
1692 "only constants or undefs allowed for "
1693 "SpvOpConstantComposite");
1694 /* to make it easier, just insert a NULL constant for now */
1695 elems[i] = vtn_null_constant(b, val->type);
1696 }
1697 }
1698
1699 switch (val->type->base_type) {
1700 case vtn_base_type_vector: {
1701 assert(glsl_type_is_vector(val->type->type));
1702 for (unsigned i = 0; i < elem_count; i++)
1703 val->constant->values[i] = elems[i]->values[0];
1704 break;
1705 }
1706
1707 case vtn_base_type_matrix:
1708 case vtn_base_type_struct:
1709 case vtn_base_type_array:
1710 ralloc_steal(val->constant, elems);
1711 val->constant->num_elements = elem_count;
1712 val->constant->elements = elems;
1713 break;
1714
1715 default:
1716 vtn_fail("Result type of %s must be a composite type",
1717 spirv_op_to_string(opcode));
1718 }
1719 break;
1720 }
1721
1722 case SpvOpSpecConstantOp: {
1723 SpvOp opcode = get_specialization(b, val, w[3]);
1724 switch (opcode) {
1725 case SpvOpVectorShuffle: {
1726 struct vtn_value *v0 = &b->values[w[4]];
1727 struct vtn_value *v1 = &b->values[w[5]];
1728
1729 vtn_assert(v0->value_type == vtn_value_type_constant ||
1730 v0->value_type == vtn_value_type_undef);
1731 vtn_assert(v1->value_type == vtn_value_type_constant ||
1732 v1->value_type == vtn_value_type_undef);
1733
1734 unsigned len0 = glsl_get_vector_elements(v0->type->type);
1735 unsigned len1 = glsl_get_vector_elements(v1->type->type);
1736
1737 vtn_assert(len0 + len1 < 16);
1738
1739 unsigned bit_size = glsl_get_bit_size(val->type->type);
1740 unsigned bit_size0 = glsl_get_bit_size(v0->type->type);
1741 unsigned bit_size1 = glsl_get_bit_size(v1->type->type);
1742
1743 vtn_assert(bit_size == bit_size0 && bit_size == bit_size1);
1744 (void)bit_size0; (void)bit_size1;
1745
1746 nir_const_value undef = { .u64 = 0xdeadbeefdeadbeef };
1747 nir_const_value combined[NIR_MAX_VEC_COMPONENTS * 2];
1748
1749 if (v0->value_type == vtn_value_type_constant) {
1750 for (unsigned i = 0; i < len0; i++)
1751 combined[i] = v0->constant->values[i];
1752 }
1753 if (v1->value_type == vtn_value_type_constant) {
1754 for (unsigned i = 0; i < len1; i++)
1755 combined[len0 + i] = v1->constant->values[i];
1756 }
1757
1758 for (unsigned i = 0, j = 0; i < count - 6; i++, j++) {
1759 uint32_t comp = w[i + 6];
1760 if (comp == (uint32_t)-1) {
1761 /* If component is not used, set the value to a known constant
1762 * to detect if it is wrongly used.
1763 */
1764 val->constant->values[j] = undef;
1765 } else {
1766 vtn_fail_if(comp >= len0 + len1,
1767 "All Component literals must either be FFFFFFFF "
1768 "or in [0, N - 1] (inclusive).");
1769 val->constant->values[j] = combined[comp];
1770 }
1771 }
1772 break;
1773 }
1774
1775 case SpvOpCompositeExtract:
1776 case SpvOpCompositeInsert: {
1777 struct vtn_value *comp;
1778 unsigned deref_start;
1779 struct nir_constant **c;
1780 if (opcode == SpvOpCompositeExtract) {
1781 comp = vtn_value(b, w[4], vtn_value_type_constant);
1782 deref_start = 5;
1783 c = &comp->constant;
1784 } else {
1785 comp = vtn_value(b, w[5], vtn_value_type_constant);
1786 deref_start = 6;
1787 val->constant = nir_constant_clone(comp->constant,
1788 (nir_variable *)b);
1789 c = &val->constant;
1790 }
1791
1792 int elem = -1;
1793 const struct vtn_type *type = comp->type;
1794 for (unsigned i = deref_start; i < count; i++) {
1795 vtn_fail_if(w[i] > type->length,
1796 "%uth index of %s is %u but the type has only "
1797 "%u elements", i - deref_start,
1798 spirv_op_to_string(opcode), w[i], type->length);
1799
1800 switch (type->base_type) {
1801 case vtn_base_type_vector:
1802 elem = w[i];
1803 type = type->array_element;
1804 break;
1805
1806 case vtn_base_type_matrix:
1807 case vtn_base_type_array:
1808 c = &(*c)->elements[w[i]];
1809 type = type->array_element;
1810 break;
1811
1812 case vtn_base_type_struct:
1813 c = &(*c)->elements[w[i]];
1814 type = type->members[w[i]];
1815 break;
1816
1817 default:
1818 vtn_fail("%s must only index into composite types",
1819 spirv_op_to_string(opcode));
1820 }
1821 }
1822
1823 if (opcode == SpvOpCompositeExtract) {
1824 if (elem == -1) {
1825 val->constant = *c;
1826 } else {
1827 unsigned num_components = type->length;
1828 for (unsigned i = 0; i < num_components; i++)
1829 val->constant->values[i] = (*c)->values[elem + i];
1830 }
1831 } else {
1832 struct vtn_value *insert =
1833 vtn_value(b, w[4], vtn_value_type_constant);
1834 vtn_assert(insert->type == type);
1835 if (elem == -1) {
1836 *c = insert->constant;
1837 } else {
1838 unsigned num_components = type->length;
1839 for (unsigned i = 0; i < num_components; i++)
1840 (*c)->values[elem + i] = insert->constant->values[i];
1841 }
1842 }
1843 break;
1844 }
1845
1846 default: {
1847 bool swap;
1848 nir_alu_type dst_alu_type = nir_get_nir_type_for_glsl_type(val->type->type);
1849 nir_alu_type src_alu_type = dst_alu_type;
1850 unsigned num_components = glsl_get_vector_elements(val->type->type);
1851 unsigned bit_size;
1852
1853 vtn_assert(count <= 7);
1854
1855 switch (opcode) {
1856 case SpvOpSConvert:
1857 case SpvOpFConvert:
1858 case SpvOpUConvert:
1859 /* We have a source in a conversion */
1860 src_alu_type =
1861 nir_get_nir_type_for_glsl_type(
1862 vtn_value(b, w[4], vtn_value_type_constant)->type->type);
1863 /* We use the bitsize of the conversion source to evaluate the opcode later */
1864 bit_size = glsl_get_bit_size(
1865 vtn_value(b, w[4], vtn_value_type_constant)->type->type);
1866 break;
1867 default:
1868 bit_size = glsl_get_bit_size(val->type->type);
1869 };
1870
1871 nir_op op = vtn_nir_alu_op_for_spirv_opcode(b, opcode, &swap,
1872 nir_alu_type_get_type_size(src_alu_type),
1873 nir_alu_type_get_type_size(dst_alu_type));
1874 nir_const_value src[3][NIR_MAX_VEC_COMPONENTS];
1875
1876 for (unsigned i = 0; i < count - 4; i++) {
1877 struct vtn_value *src_val =
1878 vtn_value(b, w[4 + i], vtn_value_type_constant);
1879
1880 /* If this is an unsized source, pull the bit size from the
1881 * source; otherwise, we'll use the bit size from the destination.
1882 */
1883 if (!nir_alu_type_get_type_size(nir_op_infos[op].input_types[i]))
1884 bit_size = glsl_get_bit_size(src_val->type->type);
1885
1886 unsigned src_comps = nir_op_infos[op].input_sizes[i] ?
1887 nir_op_infos[op].input_sizes[i] :
1888 num_components;
1889
1890 unsigned j = swap ? 1 - i : i;
1891 for (unsigned c = 0; c < src_comps; c++)
1892 src[j][c] = src_val->constant->values[c];
1893 }
1894
1895 /* fix up fixed size sources */
1896 switch (op) {
1897 case nir_op_ishl:
1898 case nir_op_ishr:
1899 case nir_op_ushr: {
1900 if (bit_size == 32)
1901 break;
1902 for (unsigned i = 0; i < num_components; ++i) {
1903 switch (bit_size) {
1904 case 64: src[1][i].u32 = src[1][i].u64; break;
1905 case 16: src[1][i].u32 = src[1][i].u16; break;
1906 case 8: src[1][i].u32 = src[1][i].u8; break;
1907 }
1908 }
1909 break;
1910 }
1911 default:
1912 break;
1913 }
1914
1915 nir_const_value *srcs[3] = {
1916 src[0], src[1], src[2],
1917 };
1918 nir_eval_const_opcode(op, val->constant->values,
1919 num_components, bit_size, srcs,
1920 b->shader->info.float_controls_execution_mode);
1921 break;
1922 } /* default */
1923 }
1924 break;
1925 }
1926
1927 case SpvOpConstantNull:
1928 val->constant = vtn_null_constant(b, val->type);
1929 break;
1930
1931 case SpvOpConstantSampler:
1932 vtn_fail("OpConstantSampler requires Kernel Capability");
1933 break;
1934
1935 default:
1936 vtn_fail_with_opcode("Unhandled opcode", opcode);
1937 }
1938
1939 /* Now that we have the value, update the workgroup size if needed */
1940 vtn_foreach_decoration(b, val, handle_workgroup_size_decoration_cb, NULL);
1941 }
1942
1943 SpvMemorySemanticsMask
1944 vtn_storage_class_to_memory_semantics(SpvStorageClass sc)
1945 {
1946 switch (sc) {
1947 case SpvStorageClassStorageBuffer:
1948 case SpvStorageClassPhysicalStorageBuffer:
1949 return SpvMemorySemanticsUniformMemoryMask;
1950 case SpvStorageClassWorkgroup:
1951 return SpvMemorySemanticsWorkgroupMemoryMask;
1952 default:
1953 return SpvMemorySemanticsMaskNone;
1954 }
1955 }
1956
1957 static void
1958 vtn_split_barrier_semantics(struct vtn_builder *b,
1959 SpvMemorySemanticsMask semantics,
1960 SpvMemorySemanticsMask *before,
1961 SpvMemorySemanticsMask *after)
1962 {
1963 /* For memory semantics embedded in operations, we split them into up to
1964 * two barriers, to be added before and after the operation. This is less
1965 * strict than if we propagated until the final backend stage, but still
1966 * result in correct execution.
1967 *
1968 * A further improvement could be pipe this information (and use!) into the
1969 * next compiler layers, at the expense of making the handling of barriers
1970 * more complicated.
1971 */
1972
1973 *before = SpvMemorySemanticsMaskNone;
1974 *after = SpvMemorySemanticsMaskNone;
1975
1976 SpvMemorySemanticsMask order_semantics =
1977 semantics & (SpvMemorySemanticsAcquireMask |
1978 SpvMemorySemanticsReleaseMask |
1979 SpvMemorySemanticsAcquireReleaseMask |
1980 SpvMemorySemanticsSequentiallyConsistentMask);
1981
1982 if (util_bitcount(order_semantics) > 1) {
1983 /* Old GLSLang versions incorrectly set all the ordering bits. This was
1984 * fixed in c51287d744fb6e7e9ccc09f6f8451e6c64b1dad6 of glslang repo,
1985 * and it is in GLSLang since revision "SPIRV99.1321" (from Jul-2016).
1986 */
1987 vtn_warn("Multiple memory ordering semantics specified, "
1988 "assuming AcquireRelease.");
1989 order_semantics = SpvMemorySemanticsAcquireReleaseMask;
1990 }
1991
1992 const SpvMemorySemanticsMask av_vis_semantics =
1993 semantics & (SpvMemorySemanticsMakeAvailableMask |
1994 SpvMemorySemanticsMakeVisibleMask);
1995
1996 const SpvMemorySemanticsMask storage_semantics =
1997 semantics & (SpvMemorySemanticsUniformMemoryMask |
1998 SpvMemorySemanticsSubgroupMemoryMask |
1999 SpvMemorySemanticsWorkgroupMemoryMask |
2000 SpvMemorySemanticsCrossWorkgroupMemoryMask |
2001 SpvMemorySemanticsAtomicCounterMemoryMask |
2002 SpvMemorySemanticsImageMemoryMask |
2003 SpvMemorySemanticsOutputMemoryMask);
2004
2005 const SpvMemorySemanticsMask other_semantics =
2006 semantics & ~(order_semantics | av_vis_semantics | storage_semantics);
2007
2008 if (other_semantics)
2009 vtn_warn("Ignoring unhandled memory semantics: %u\n", other_semantics);
2010
2011 /* SequentiallyConsistent is treated as AcquireRelease. */
2012
2013 /* The RELEASE barrier happens BEFORE the operation, and it is usually
2014 * associated with a Store. All the write operations with a matching
2015 * semantics will not be reordered after the Store.
2016 */
2017 if (order_semantics & (SpvMemorySemanticsReleaseMask |
2018 SpvMemorySemanticsAcquireReleaseMask |
2019 SpvMemorySemanticsSequentiallyConsistentMask)) {
2020 *before |= SpvMemorySemanticsReleaseMask | storage_semantics;
2021 }
2022
2023 /* The ACQUIRE barrier happens AFTER the operation, and it is usually
2024 * associated with a Load. All the operations with a matching semantics
2025 * will not be reordered before the Load.
2026 */
2027 if (order_semantics & (SpvMemorySemanticsAcquireMask |
2028 SpvMemorySemanticsAcquireReleaseMask |
2029 SpvMemorySemanticsSequentiallyConsistentMask)) {
2030 *after |= SpvMemorySemanticsAcquireMask | storage_semantics;
2031 }
2032
2033 if (av_vis_semantics & SpvMemorySemanticsMakeVisibleMask)
2034 *before |= SpvMemorySemanticsMakeVisibleMask | storage_semantics;
2035
2036 if (av_vis_semantics & SpvMemorySemanticsMakeAvailableMask)
2037 *after |= SpvMemorySemanticsMakeAvailableMask | storage_semantics;
2038 }
2039
2040 static void
2041 vtn_emit_scoped_memory_barrier(struct vtn_builder *b, SpvScope scope,
2042 SpvMemorySemanticsMask semantics)
2043 {
2044 nir_memory_semantics nir_semantics = 0;
2045
2046 SpvMemorySemanticsMask order_semantics =
2047 semantics & (SpvMemorySemanticsAcquireMask |
2048 SpvMemorySemanticsReleaseMask |
2049 SpvMemorySemanticsAcquireReleaseMask |
2050 SpvMemorySemanticsSequentiallyConsistentMask);
2051
2052 if (util_bitcount(order_semantics) > 1) {
2053 /* Old GLSLang versions incorrectly set all the ordering bits. This was
2054 * fixed in c51287d744fb6e7e9ccc09f6f8451e6c64b1dad6 of glslang repo,
2055 * and it is in GLSLang since revision "SPIRV99.1321" (from Jul-2016).
2056 */
2057 vtn_warn("Multiple memory ordering semantics bits specified, "
2058 "assuming AcquireRelease.");
2059 order_semantics = SpvMemorySemanticsAcquireReleaseMask;
2060 }
2061
2062 switch (order_semantics) {
2063 case 0:
2064 /* Not an ordering barrier. */
2065 break;
2066
2067 case SpvMemorySemanticsAcquireMask:
2068 nir_semantics = NIR_MEMORY_ACQUIRE;
2069 break;
2070
2071 case SpvMemorySemanticsReleaseMask:
2072 nir_semantics = NIR_MEMORY_RELEASE;
2073 break;
2074
2075 case SpvMemorySemanticsSequentiallyConsistentMask:
2076 /* Fall through. Treated as AcquireRelease in Vulkan. */
2077 case SpvMemorySemanticsAcquireReleaseMask:
2078 nir_semantics = NIR_MEMORY_ACQUIRE | NIR_MEMORY_RELEASE;
2079 break;
2080
2081 default:
2082 unreachable("Invalid memory order semantics");
2083 }
2084
2085 if (semantics & SpvMemorySemanticsMakeAvailableMask) {
2086 vtn_fail_if(!b->options->caps.vk_memory_model,
2087 "To use MakeAvailable memory semantics the VulkanMemoryModel "
2088 "capability must be declared.");
2089 nir_semantics |= NIR_MEMORY_MAKE_AVAILABLE;
2090 }
2091
2092 if (semantics & SpvMemorySemanticsMakeVisibleMask) {
2093 vtn_fail_if(!b->options->caps.vk_memory_model,
2094 "To use MakeVisible memory semantics the VulkanMemoryModel "
2095 "capability must be declared.");
2096 nir_semantics |= NIR_MEMORY_MAKE_VISIBLE;
2097 }
2098
2099 /* Vulkan Environment for SPIR-V says "SubgroupMemory, CrossWorkgroupMemory,
2100 * and AtomicCounterMemory are ignored".
2101 */
2102 semantics &= ~(SpvMemorySemanticsSubgroupMemoryMask |
2103 SpvMemorySemanticsCrossWorkgroupMemoryMask |
2104 SpvMemorySemanticsAtomicCounterMemoryMask);
2105
2106 /* TODO: Consider adding nir_var_mem_image mode to NIR so it can be used
2107 * for SpvMemorySemanticsImageMemoryMask.
2108 */
2109
2110 nir_variable_mode modes = 0;
2111 if (semantics & (SpvMemorySemanticsUniformMemoryMask |
2112 SpvMemorySemanticsImageMemoryMask)) {
2113 modes |= nir_var_uniform |
2114 nir_var_mem_ubo |
2115 nir_var_mem_ssbo |
2116 nir_var_mem_global;
2117 }
2118 if (semantics & SpvMemorySemanticsWorkgroupMemoryMask)
2119 modes |= nir_var_mem_shared;
2120 if (semantics & SpvMemorySemanticsOutputMemoryMask) {
2121 modes |= nir_var_shader_out;
2122 }
2123
2124 /* No barrier to add. */
2125 if (nir_semantics == 0 || modes == 0)
2126 return;
2127
2128 nir_scope nir_scope;
2129 switch (scope) {
2130 case SpvScopeDevice:
2131 vtn_fail_if(b->options->caps.vk_memory_model &&
2132 !b->options->caps.vk_memory_model_device_scope,
2133 "If the Vulkan memory model is declared and any instruction "
2134 "uses Device scope, the VulkanMemoryModelDeviceScope "
2135 "capability must be declared.");
2136 nir_scope = NIR_SCOPE_DEVICE;
2137 break;
2138
2139 case SpvScopeQueueFamily:
2140 vtn_fail_if(!b->options->caps.vk_memory_model,
2141 "To use Queue Family scope, the VulkanMemoryModel capability "
2142 "must be declared.");
2143 nir_scope = NIR_SCOPE_QUEUE_FAMILY;
2144 break;
2145
2146 case SpvScopeWorkgroup:
2147 nir_scope = NIR_SCOPE_WORKGROUP;
2148 break;
2149
2150 case SpvScopeSubgroup:
2151 nir_scope = NIR_SCOPE_SUBGROUP;
2152 break;
2153
2154 case SpvScopeInvocation:
2155 nir_scope = NIR_SCOPE_INVOCATION;
2156 break;
2157
2158 default:
2159 vtn_fail("Invalid memory scope");
2160 }
2161
2162 nir_intrinsic_instr *intrin =
2163 nir_intrinsic_instr_create(b->shader, nir_intrinsic_scoped_memory_barrier);
2164 nir_intrinsic_set_memory_semantics(intrin, nir_semantics);
2165
2166 nir_intrinsic_set_memory_modes(intrin, modes);
2167 nir_intrinsic_set_memory_scope(intrin, nir_scope);
2168 nir_builder_instr_insert(&b->nb, &intrin->instr);
2169 }
2170
2171 struct vtn_ssa_value *
2172 vtn_create_ssa_value(struct vtn_builder *b, const struct glsl_type *type)
2173 {
2174 struct vtn_ssa_value *val = rzalloc(b, struct vtn_ssa_value);
2175 val->type = type;
2176
2177 if (!glsl_type_is_vector_or_scalar(type)) {
2178 unsigned elems = glsl_get_length(type);
2179 val->elems = ralloc_array(b, struct vtn_ssa_value *, elems);
2180 for (unsigned i = 0; i < elems; i++) {
2181 const struct glsl_type *child_type;
2182
2183 switch (glsl_get_base_type(type)) {
2184 case GLSL_TYPE_INT:
2185 case GLSL_TYPE_UINT:
2186 case GLSL_TYPE_INT16:
2187 case GLSL_TYPE_UINT16:
2188 case GLSL_TYPE_UINT8:
2189 case GLSL_TYPE_INT8:
2190 case GLSL_TYPE_INT64:
2191 case GLSL_TYPE_UINT64:
2192 case GLSL_TYPE_BOOL:
2193 case GLSL_TYPE_FLOAT:
2194 case GLSL_TYPE_FLOAT16:
2195 case GLSL_TYPE_DOUBLE:
2196 child_type = glsl_get_column_type(type);
2197 break;
2198 case GLSL_TYPE_ARRAY:
2199 child_type = glsl_get_array_element(type);
2200 break;
2201 case GLSL_TYPE_STRUCT:
2202 case GLSL_TYPE_INTERFACE:
2203 child_type = glsl_get_struct_field(type, i);
2204 break;
2205 default:
2206 vtn_fail("unkown base type");
2207 }
2208
2209 val->elems[i] = vtn_create_ssa_value(b, child_type);
2210 }
2211 }
2212
2213 return val;
2214 }
2215
2216 static nir_tex_src
2217 vtn_tex_src(struct vtn_builder *b, unsigned index, nir_tex_src_type type)
2218 {
2219 nir_tex_src src;
2220 src.src = nir_src_for_ssa(vtn_ssa_value(b, index)->def);
2221 src.src_type = type;
2222 return src;
2223 }
2224
2225 static uint32_t
2226 image_operand_arg(struct vtn_builder *b, const uint32_t *w, uint32_t count,
2227 uint32_t mask_idx, SpvImageOperandsMask op)
2228 {
2229 static const SpvImageOperandsMask ops_with_arg =
2230 SpvImageOperandsBiasMask |
2231 SpvImageOperandsLodMask |
2232 SpvImageOperandsGradMask |
2233 SpvImageOperandsConstOffsetMask |
2234 SpvImageOperandsOffsetMask |
2235 SpvImageOperandsConstOffsetsMask |
2236 SpvImageOperandsSampleMask |
2237 SpvImageOperandsMinLodMask |
2238 SpvImageOperandsMakeTexelAvailableMask |
2239 SpvImageOperandsMakeTexelVisibleMask;
2240
2241 assert(util_bitcount(op) == 1);
2242 assert(w[mask_idx] & op);
2243 assert(op & ops_with_arg);
2244
2245 uint32_t idx = util_bitcount(w[mask_idx] & (op - 1) & ops_with_arg) + 1;
2246
2247 /* Adjust indices for operands with two arguments. */
2248 static const SpvImageOperandsMask ops_with_two_args =
2249 SpvImageOperandsGradMask;
2250 idx += util_bitcount(w[mask_idx] & (op - 1) & ops_with_two_args);
2251
2252 idx += mask_idx;
2253
2254 vtn_fail_if(idx + (op & ops_with_two_args ? 1 : 0) >= count,
2255 "Image op claims to have %s but does not enough "
2256 "following operands", spirv_imageoperands_to_string(op));
2257
2258 return idx;
2259 }
2260
2261 static void
2262 vtn_handle_texture(struct vtn_builder *b, SpvOp opcode,
2263 const uint32_t *w, unsigned count)
2264 {
2265 if (opcode == SpvOpSampledImage) {
2266 struct vtn_value *val =
2267 vtn_push_value(b, w[2], vtn_value_type_sampled_image);
2268 val->sampled_image = ralloc(b, struct vtn_sampled_image);
2269 val->sampled_image->image =
2270 vtn_value(b, w[3], vtn_value_type_pointer)->pointer;
2271 val->sampled_image->sampler =
2272 vtn_value(b, w[4], vtn_value_type_pointer)->pointer;
2273 return;
2274 } else if (opcode == SpvOpImage) {
2275 struct vtn_value *src_val = vtn_untyped_value(b, w[3]);
2276 if (src_val->value_type == vtn_value_type_sampled_image) {
2277 vtn_push_value_pointer(b, w[2], src_val->sampled_image->image);
2278 } else {
2279 vtn_assert(src_val->value_type == vtn_value_type_pointer);
2280 vtn_push_value_pointer(b, w[2], src_val->pointer);
2281 }
2282 return;
2283 }
2284
2285 struct vtn_type *ret_type = vtn_value(b, w[1], vtn_value_type_type)->type;
2286
2287 struct vtn_pointer *image = NULL, *sampler = NULL;
2288 struct vtn_value *sampled_val = vtn_untyped_value(b, w[3]);
2289 if (sampled_val->value_type == vtn_value_type_sampled_image) {
2290 image = sampled_val->sampled_image->image;
2291 sampler = sampled_val->sampled_image->sampler;
2292 } else {
2293 vtn_assert(sampled_val->value_type == vtn_value_type_pointer);
2294 image = sampled_val->pointer;
2295 }
2296
2297 nir_deref_instr *image_deref = vtn_pointer_to_deref(b, image);
2298 nir_deref_instr *sampler_deref =
2299 sampler ? vtn_pointer_to_deref(b, sampler) : NULL;
2300
2301 const struct glsl_type *image_type = sampled_val->type->type;
2302 const enum glsl_sampler_dim sampler_dim = glsl_get_sampler_dim(image_type);
2303 const bool is_array = glsl_sampler_type_is_array(image_type);
2304 nir_alu_type dest_type = nir_type_invalid;
2305
2306 /* Figure out the base texture operation */
2307 nir_texop texop;
2308 switch (opcode) {
2309 case SpvOpImageSampleImplicitLod:
2310 case SpvOpImageSampleDrefImplicitLod:
2311 case SpvOpImageSampleProjImplicitLod:
2312 case SpvOpImageSampleProjDrefImplicitLod:
2313 texop = nir_texop_tex;
2314 break;
2315
2316 case SpvOpImageSampleExplicitLod:
2317 case SpvOpImageSampleDrefExplicitLod:
2318 case SpvOpImageSampleProjExplicitLod:
2319 case SpvOpImageSampleProjDrefExplicitLod:
2320 texop = nir_texop_txl;
2321 break;
2322
2323 case SpvOpImageFetch:
2324 if (sampler_dim == GLSL_SAMPLER_DIM_MS) {
2325 texop = nir_texop_txf_ms;
2326 } else {
2327 texop = nir_texop_txf;
2328 }
2329 break;
2330
2331 case SpvOpImageGather:
2332 case SpvOpImageDrefGather:
2333 texop = nir_texop_tg4;
2334 break;
2335
2336 case SpvOpImageQuerySizeLod:
2337 case SpvOpImageQuerySize:
2338 texop = nir_texop_txs;
2339 dest_type = nir_type_int;
2340 break;
2341
2342 case SpvOpImageQueryLod:
2343 texop = nir_texop_lod;
2344 dest_type = nir_type_float;
2345 break;
2346
2347 case SpvOpImageQueryLevels:
2348 texop = nir_texop_query_levels;
2349 dest_type = nir_type_int;
2350 break;
2351
2352 case SpvOpImageQuerySamples:
2353 texop = nir_texop_texture_samples;
2354 dest_type = nir_type_int;
2355 break;
2356
2357 case SpvOpFragmentFetchAMD:
2358 texop = nir_texop_fragment_fetch;
2359 break;
2360
2361 case SpvOpFragmentMaskFetchAMD:
2362 texop = nir_texop_fragment_mask_fetch;
2363 break;
2364
2365 default:
2366 vtn_fail_with_opcode("Unhandled opcode", opcode);
2367 }
2368
2369 nir_tex_src srcs[10]; /* 10 should be enough */
2370 nir_tex_src *p = srcs;
2371
2372 p->src = nir_src_for_ssa(&image_deref->dest.ssa);
2373 p->src_type = nir_tex_src_texture_deref;
2374 p++;
2375
2376 switch (texop) {
2377 case nir_texop_tex:
2378 case nir_texop_txb:
2379 case nir_texop_txl:
2380 case nir_texop_txd:
2381 case nir_texop_tg4:
2382 case nir_texop_lod:
2383 vtn_fail_if(sampler == NULL,
2384 "%s requires an image of type OpTypeSampledImage",
2385 spirv_op_to_string(opcode));
2386 p->src = nir_src_for_ssa(&sampler_deref->dest.ssa);
2387 p->src_type = nir_tex_src_sampler_deref;
2388 p++;
2389 break;
2390 case nir_texop_txf:
2391 case nir_texop_txf_ms:
2392 case nir_texop_txs:
2393 case nir_texop_query_levels:
2394 case nir_texop_texture_samples:
2395 case nir_texop_samples_identical:
2396 case nir_texop_fragment_fetch:
2397 case nir_texop_fragment_mask_fetch:
2398 /* These don't */
2399 break;
2400 case nir_texop_txf_ms_fb:
2401 vtn_fail("unexpected nir_texop_txf_ms_fb");
2402 break;
2403 case nir_texop_txf_ms_mcs:
2404 vtn_fail("unexpected nir_texop_txf_ms_mcs");
2405 case nir_texop_tex_prefetch:
2406 vtn_fail("unexpected nir_texop_tex_prefetch");
2407 }
2408
2409 unsigned idx = 4;
2410
2411 struct nir_ssa_def *coord;
2412 unsigned coord_components;
2413 switch (opcode) {
2414 case SpvOpImageSampleImplicitLod:
2415 case SpvOpImageSampleExplicitLod:
2416 case SpvOpImageSampleDrefImplicitLod:
2417 case SpvOpImageSampleDrefExplicitLod:
2418 case SpvOpImageSampleProjImplicitLod:
2419 case SpvOpImageSampleProjExplicitLod:
2420 case SpvOpImageSampleProjDrefImplicitLod:
2421 case SpvOpImageSampleProjDrefExplicitLod:
2422 case SpvOpImageFetch:
2423 case SpvOpImageGather:
2424 case SpvOpImageDrefGather:
2425 case SpvOpImageQueryLod:
2426 case SpvOpFragmentFetchAMD:
2427 case SpvOpFragmentMaskFetchAMD: {
2428 /* All these types have the coordinate as their first real argument */
2429 switch (sampler_dim) {
2430 case GLSL_SAMPLER_DIM_1D:
2431 case GLSL_SAMPLER_DIM_BUF:
2432 coord_components = 1;
2433 break;
2434 case GLSL_SAMPLER_DIM_2D:
2435 case GLSL_SAMPLER_DIM_RECT:
2436 case GLSL_SAMPLER_DIM_MS:
2437 case GLSL_SAMPLER_DIM_SUBPASS_MS:
2438 coord_components = 2;
2439 break;
2440 case GLSL_SAMPLER_DIM_3D:
2441 case GLSL_SAMPLER_DIM_CUBE:
2442 coord_components = 3;
2443 break;
2444 default:
2445 vtn_fail("Invalid sampler type");
2446 }
2447
2448 if (is_array && texop != nir_texop_lod)
2449 coord_components++;
2450
2451 coord = vtn_ssa_value(b, w[idx++])->def;
2452 p->src = nir_src_for_ssa(nir_channels(&b->nb, coord,
2453 (1 << coord_components) - 1));
2454 p->src_type = nir_tex_src_coord;
2455 p++;
2456 break;
2457 }
2458
2459 default:
2460 coord = NULL;
2461 coord_components = 0;
2462 break;
2463 }
2464
2465 switch (opcode) {
2466 case SpvOpImageSampleProjImplicitLod:
2467 case SpvOpImageSampleProjExplicitLod:
2468 case SpvOpImageSampleProjDrefImplicitLod:
2469 case SpvOpImageSampleProjDrefExplicitLod:
2470 /* These have the projector as the last coordinate component */
2471 p->src = nir_src_for_ssa(nir_channel(&b->nb, coord, coord_components));
2472 p->src_type = nir_tex_src_projector;
2473 p++;
2474 break;
2475
2476 default:
2477 break;
2478 }
2479
2480 bool is_shadow = false;
2481 unsigned gather_component = 0;
2482 switch (opcode) {
2483 case SpvOpImageSampleDrefImplicitLod:
2484 case SpvOpImageSampleDrefExplicitLod:
2485 case SpvOpImageSampleProjDrefImplicitLod:
2486 case SpvOpImageSampleProjDrefExplicitLod:
2487 case SpvOpImageDrefGather:
2488 /* These all have an explicit depth value as their next source */
2489 is_shadow = true;
2490 (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_comparator);
2491 break;
2492
2493 case SpvOpImageGather:
2494 /* This has a component as its next source */
2495 gather_component = vtn_constant_uint(b, w[idx++]);
2496 break;
2497
2498 default:
2499 break;
2500 }
2501
2502 /* For OpImageQuerySizeLod, we always have an LOD */
2503 if (opcode == SpvOpImageQuerySizeLod)
2504 (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_lod);
2505
2506 /* For OpFragmentFetchAMD, we always have a multisample index */
2507 if (opcode == SpvOpFragmentFetchAMD)
2508 (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_ms_index);
2509
2510 /* Now we need to handle some number of optional arguments */
2511 struct vtn_value *gather_offsets = NULL;
2512 if (idx < count) {
2513 uint32_t operands = w[idx];
2514
2515 if (operands & SpvImageOperandsBiasMask) {
2516 vtn_assert(texop == nir_texop_tex);
2517 texop = nir_texop_txb;
2518 uint32_t arg = image_operand_arg(b, w, count, idx,
2519 SpvImageOperandsBiasMask);
2520 (*p++) = vtn_tex_src(b, w[arg], nir_tex_src_bias);
2521 }
2522
2523 if (operands & SpvImageOperandsLodMask) {
2524 vtn_assert(texop == nir_texop_txl || texop == nir_texop_txf ||
2525 texop == nir_texop_txs);
2526 uint32_t arg = image_operand_arg(b, w, count, idx,
2527 SpvImageOperandsLodMask);
2528 (*p++) = vtn_tex_src(b, w[arg], nir_tex_src_lod);
2529 }
2530
2531 if (operands & SpvImageOperandsGradMask) {
2532 vtn_assert(texop == nir_texop_txl);
2533 texop = nir_texop_txd;
2534 uint32_t arg = image_operand_arg(b, w, count, idx,
2535 SpvImageOperandsGradMask);
2536 (*p++) = vtn_tex_src(b, w[arg], nir_tex_src_ddx);
2537 (*p++) = vtn_tex_src(b, w[arg + 1], nir_tex_src_ddy);
2538 }
2539
2540 vtn_fail_if(util_bitcount(operands & (SpvImageOperandsConstOffsetsMask |
2541 SpvImageOperandsOffsetMask |
2542 SpvImageOperandsConstOffsetMask)) > 1,
2543 "At most one of the ConstOffset, Offset, and ConstOffsets "
2544 "image operands can be used on a given instruction.");
2545
2546 if (operands & SpvImageOperandsOffsetMask) {
2547 uint32_t arg = image_operand_arg(b, w, count, idx,
2548 SpvImageOperandsOffsetMask);
2549 (*p++) = vtn_tex_src(b, w[arg], nir_tex_src_offset);
2550 }
2551
2552 if (operands & SpvImageOperandsConstOffsetMask) {
2553 uint32_t arg = image_operand_arg(b, w, count, idx,
2554 SpvImageOperandsConstOffsetMask);
2555 (*p++) = vtn_tex_src(b, w[arg], nir_tex_src_offset);
2556 }
2557
2558 if (operands & SpvImageOperandsConstOffsetsMask) {
2559 vtn_assert(texop == nir_texop_tg4);
2560 uint32_t arg = image_operand_arg(b, w, count, idx,
2561 SpvImageOperandsConstOffsetsMask);
2562 gather_offsets = vtn_value(b, w[arg], vtn_value_type_constant);
2563 }
2564
2565 if (operands & SpvImageOperandsSampleMask) {
2566 vtn_assert(texop == nir_texop_txf_ms);
2567 uint32_t arg = image_operand_arg(b, w, count, idx,
2568 SpvImageOperandsSampleMask);
2569 texop = nir_texop_txf_ms;
2570 (*p++) = vtn_tex_src(b, w[arg], nir_tex_src_ms_index);
2571 }
2572
2573 if (operands & SpvImageOperandsMinLodMask) {
2574 vtn_assert(texop == nir_texop_tex ||
2575 texop == nir_texop_txb ||
2576 texop == nir_texop_txd);
2577 uint32_t arg = image_operand_arg(b, w, count, idx,
2578 SpvImageOperandsMinLodMask);
2579 (*p++) = vtn_tex_src(b, w[arg], nir_tex_src_min_lod);
2580 }
2581 }
2582
2583 nir_tex_instr *instr = nir_tex_instr_create(b->shader, p - srcs);
2584 instr->op = texop;
2585
2586 memcpy(instr->src, srcs, instr->num_srcs * sizeof(*instr->src));
2587
2588 instr->coord_components = coord_components;
2589 instr->sampler_dim = sampler_dim;
2590 instr->is_array = is_array;
2591 instr->is_shadow = is_shadow;
2592 instr->is_new_style_shadow =
2593 is_shadow && glsl_get_components(ret_type->type) == 1;
2594 instr->component = gather_component;
2595
2596 if (image && (image->access & ACCESS_NON_UNIFORM))
2597 instr->texture_non_uniform = true;
2598
2599 if (sampler && (sampler->access & ACCESS_NON_UNIFORM))
2600 instr->sampler_non_uniform = true;
2601
2602 /* for non-query ops, get dest_type from sampler type */
2603 if (dest_type == nir_type_invalid) {
2604 switch (glsl_get_sampler_result_type(image_type)) {
2605 case GLSL_TYPE_FLOAT: dest_type = nir_type_float; break;
2606 case GLSL_TYPE_INT: dest_type = nir_type_int; break;
2607 case GLSL_TYPE_UINT: dest_type = nir_type_uint; break;
2608 case GLSL_TYPE_BOOL: dest_type = nir_type_bool; break;
2609 default:
2610 vtn_fail("Invalid base type for sampler result");
2611 }
2612 }
2613
2614 instr->dest_type = dest_type;
2615
2616 nir_ssa_dest_init(&instr->instr, &instr->dest,
2617 nir_tex_instr_dest_size(instr), 32, NULL);
2618
2619 vtn_assert(glsl_get_vector_elements(ret_type->type) ==
2620 nir_tex_instr_dest_size(instr));
2621
2622 if (gather_offsets) {
2623 vtn_fail_if(gather_offsets->type->base_type != vtn_base_type_array ||
2624 gather_offsets->type->length != 4,
2625 "ConstOffsets must be an array of size four of vectors "
2626 "of two integer components");
2627
2628 struct vtn_type *vec_type = gather_offsets->type->array_element;
2629 vtn_fail_if(vec_type->base_type != vtn_base_type_vector ||
2630 vec_type->length != 2 ||
2631 !glsl_type_is_integer(vec_type->type),
2632 "ConstOffsets must be an array of size four of vectors "
2633 "of two integer components");
2634
2635 unsigned bit_size = glsl_get_bit_size(vec_type->type);
2636 for (uint32_t i = 0; i < 4; i++) {
2637 const nir_const_value *cvec =
2638 gather_offsets->constant->elements[i]->values;
2639 for (uint32_t j = 0; j < 2; j++) {
2640 switch (bit_size) {
2641 case 8: instr->tg4_offsets[i][j] = cvec[j].i8; break;
2642 case 16: instr->tg4_offsets[i][j] = cvec[j].i16; break;
2643 case 32: instr->tg4_offsets[i][j] = cvec[j].i32; break;
2644 case 64: instr->tg4_offsets[i][j] = cvec[j].i64; break;
2645 default:
2646 vtn_fail("Unsupported bit size: %u", bit_size);
2647 }
2648 }
2649 }
2650 }
2651
2652 struct vtn_ssa_value *ssa = vtn_create_ssa_value(b, ret_type->type);
2653 ssa->def = &instr->dest.ssa;
2654 vtn_push_ssa(b, w[2], ret_type, ssa);
2655
2656 nir_builder_instr_insert(&b->nb, &instr->instr);
2657 }
2658
2659 static void
2660 fill_common_atomic_sources(struct vtn_builder *b, SpvOp opcode,
2661 const uint32_t *w, nir_src *src)
2662 {
2663 switch (opcode) {
2664 case SpvOpAtomicIIncrement:
2665 src[0] = nir_src_for_ssa(nir_imm_int(&b->nb, 1));
2666 break;
2667
2668 case SpvOpAtomicIDecrement:
2669 src[0] = nir_src_for_ssa(nir_imm_int(&b->nb, -1));
2670 break;
2671
2672 case SpvOpAtomicISub:
2673 src[0] =
2674 nir_src_for_ssa(nir_ineg(&b->nb, vtn_ssa_value(b, w[6])->def));
2675 break;
2676
2677 case SpvOpAtomicCompareExchange:
2678 case SpvOpAtomicCompareExchangeWeak:
2679 src[0] = nir_src_for_ssa(vtn_ssa_value(b, w[8])->def);
2680 src[1] = nir_src_for_ssa(vtn_ssa_value(b, w[7])->def);
2681 break;
2682
2683 case SpvOpAtomicExchange:
2684 case SpvOpAtomicIAdd:
2685 case SpvOpAtomicSMin:
2686 case SpvOpAtomicUMin:
2687 case SpvOpAtomicSMax:
2688 case SpvOpAtomicUMax:
2689 case SpvOpAtomicAnd:
2690 case SpvOpAtomicOr:
2691 case SpvOpAtomicXor:
2692 src[0] = nir_src_for_ssa(vtn_ssa_value(b, w[6])->def);
2693 break;
2694
2695 default:
2696 vtn_fail_with_opcode("Invalid SPIR-V atomic", opcode);
2697 }
2698 }
2699
2700 static nir_ssa_def *
2701 get_image_coord(struct vtn_builder *b, uint32_t value)
2702 {
2703 struct vtn_ssa_value *coord = vtn_ssa_value(b, value);
2704
2705 /* The image_load_store intrinsics assume a 4-dim coordinate */
2706 unsigned dim = glsl_get_vector_elements(coord->type);
2707 unsigned swizzle[4];
2708 for (unsigned i = 0; i < 4; i++)
2709 swizzle[i] = MIN2(i, dim - 1);
2710
2711 return nir_swizzle(&b->nb, coord->def, swizzle, 4);
2712 }
2713
2714 static nir_ssa_def *
2715 expand_to_vec4(nir_builder *b, nir_ssa_def *value)
2716 {
2717 if (value->num_components == 4)
2718 return value;
2719
2720 unsigned swiz[4];
2721 for (unsigned i = 0; i < 4; i++)
2722 swiz[i] = i < value->num_components ? i : 0;
2723 return nir_swizzle(b, value, swiz, 4);
2724 }
2725
2726 static void
2727 vtn_handle_image(struct vtn_builder *b, SpvOp opcode,
2728 const uint32_t *w, unsigned count)
2729 {
2730 /* Just get this one out of the way */
2731 if (opcode == SpvOpImageTexelPointer) {
2732 struct vtn_value *val =
2733 vtn_push_value(b, w[2], vtn_value_type_image_pointer);
2734 val->image = ralloc(b, struct vtn_image_pointer);
2735
2736 val->image->image = vtn_value(b, w[3], vtn_value_type_pointer)->pointer;
2737 val->image->coord = get_image_coord(b, w[4]);
2738 val->image->sample = vtn_ssa_value(b, w[5])->def;
2739 val->image->lod = nir_imm_int(&b->nb, 0);
2740 return;
2741 }
2742
2743 struct vtn_image_pointer image;
2744 SpvScope scope = SpvScopeInvocation;
2745 SpvMemorySemanticsMask semantics = 0;
2746
2747 switch (opcode) {
2748 case SpvOpAtomicExchange:
2749 case SpvOpAtomicCompareExchange:
2750 case SpvOpAtomicCompareExchangeWeak:
2751 case SpvOpAtomicIIncrement:
2752 case SpvOpAtomicIDecrement:
2753 case SpvOpAtomicIAdd:
2754 case SpvOpAtomicISub:
2755 case SpvOpAtomicLoad:
2756 case SpvOpAtomicSMin:
2757 case SpvOpAtomicUMin:
2758 case SpvOpAtomicSMax:
2759 case SpvOpAtomicUMax:
2760 case SpvOpAtomicAnd:
2761 case SpvOpAtomicOr:
2762 case SpvOpAtomicXor:
2763 image = *vtn_value(b, w[3], vtn_value_type_image_pointer)->image;
2764 scope = vtn_constant_uint(b, w[4]);
2765 semantics = vtn_constant_uint(b, w[5]);
2766 break;
2767
2768 case SpvOpAtomicStore:
2769 image = *vtn_value(b, w[1], vtn_value_type_image_pointer)->image;
2770 scope = vtn_constant_uint(b, w[2]);
2771 semantics = vtn_constant_uint(b, w[3]);
2772 break;
2773
2774 case SpvOpImageQuerySize:
2775 image.image = vtn_value(b, w[3], vtn_value_type_pointer)->pointer;
2776 image.coord = NULL;
2777 image.sample = NULL;
2778 image.lod = NULL;
2779 break;
2780
2781 case SpvOpImageRead: {
2782 image.image = vtn_value(b, w[3], vtn_value_type_pointer)->pointer;
2783 image.coord = get_image_coord(b, w[4]);
2784
2785 const SpvImageOperandsMask operands =
2786 count > 5 ? w[5] : SpvImageOperandsMaskNone;
2787
2788 if (operands & SpvImageOperandsSampleMask) {
2789 uint32_t arg = image_operand_arg(b, w, count, 5,
2790 SpvImageOperandsSampleMask);
2791 image.sample = vtn_ssa_value(b, w[arg])->def;
2792 } else {
2793 image.sample = nir_ssa_undef(&b->nb, 1, 32);
2794 }
2795
2796 if (operands & SpvImageOperandsMakeTexelVisibleMask) {
2797 vtn_fail_if((operands & SpvImageOperandsNonPrivateTexelMask) == 0,
2798 "MakeTexelVisible requires NonPrivateTexel to also be set.");
2799 uint32_t arg = image_operand_arg(b, w, count, 5,
2800 SpvImageOperandsMakeTexelVisibleMask);
2801 semantics = SpvMemorySemanticsMakeVisibleMask;
2802 scope = vtn_constant_uint(b, w[arg]);
2803 }
2804
2805 if (operands & SpvImageOperandsLodMask) {
2806 uint32_t arg = image_operand_arg(b, w, count, 5,
2807 SpvImageOperandsLodMask);
2808 image.lod = vtn_ssa_value(b, w[arg])->def;
2809 } else {
2810 image.lod = nir_imm_int(&b->nb, 0);
2811 }
2812
2813 /* TODO: Volatile. */
2814
2815 break;
2816 }
2817
2818 case SpvOpImageWrite: {
2819 image.image = vtn_value(b, w[1], vtn_value_type_pointer)->pointer;
2820 image.coord = get_image_coord(b, w[2]);
2821
2822 /* texel = w[3] */
2823
2824 const SpvImageOperandsMask operands =
2825 count > 4 ? w[4] : SpvImageOperandsMaskNone;
2826
2827 if (operands & SpvImageOperandsSampleMask) {
2828 uint32_t arg = image_operand_arg(b, w, count, 4,
2829 SpvImageOperandsSampleMask);
2830 image.sample = vtn_ssa_value(b, w[arg])->def;
2831 } else {
2832 image.sample = nir_ssa_undef(&b->nb, 1, 32);
2833 }
2834
2835 if (operands & SpvImageOperandsMakeTexelAvailableMask) {
2836 vtn_fail_if((operands & SpvImageOperandsNonPrivateTexelMask) == 0,
2837 "MakeTexelAvailable requires NonPrivateTexel to also be set.");
2838 uint32_t arg = image_operand_arg(b, w, count, 4,
2839 SpvImageOperandsMakeTexelAvailableMask);
2840 semantics = SpvMemorySemanticsMakeAvailableMask;
2841 scope = vtn_constant_uint(b, w[arg]);
2842 }
2843
2844 if (operands & SpvImageOperandsLodMask) {
2845 uint32_t arg = image_operand_arg(b, w, count, 4,
2846 SpvImageOperandsLodMask);
2847 image.lod = vtn_ssa_value(b, w[arg])->def;
2848 } else {
2849 image.lod = nir_imm_int(&b->nb, 0);
2850 }
2851
2852 /* TODO: Volatile. */
2853
2854 break;
2855 }
2856
2857 default:
2858 vtn_fail_with_opcode("Invalid image opcode", opcode);
2859 }
2860
2861 nir_intrinsic_op op;
2862 switch (opcode) {
2863 #define OP(S, N) case SpvOp##S: op = nir_intrinsic_image_deref_##N; break;
2864 OP(ImageQuerySize, size)
2865 OP(ImageRead, load)
2866 OP(ImageWrite, store)
2867 OP(AtomicLoad, load)
2868 OP(AtomicStore, store)
2869 OP(AtomicExchange, atomic_exchange)
2870 OP(AtomicCompareExchange, atomic_comp_swap)
2871 OP(AtomicCompareExchangeWeak, atomic_comp_swap)
2872 OP(AtomicIIncrement, atomic_add)
2873 OP(AtomicIDecrement, atomic_add)
2874 OP(AtomicIAdd, atomic_add)
2875 OP(AtomicISub, atomic_add)
2876 OP(AtomicSMin, atomic_imin)
2877 OP(AtomicUMin, atomic_umin)
2878 OP(AtomicSMax, atomic_imax)
2879 OP(AtomicUMax, atomic_umax)
2880 OP(AtomicAnd, atomic_and)
2881 OP(AtomicOr, atomic_or)
2882 OP(AtomicXor, atomic_xor)
2883 #undef OP
2884 default:
2885 vtn_fail_with_opcode("Invalid image opcode", opcode);
2886 }
2887
2888 nir_intrinsic_instr *intrin = nir_intrinsic_instr_create(b->shader, op);
2889
2890 nir_deref_instr *image_deref = vtn_pointer_to_deref(b, image.image);
2891 intrin->src[0] = nir_src_for_ssa(&image_deref->dest.ssa);
2892
2893 /* ImageQuerySize doesn't take any extra parameters */
2894 if (opcode != SpvOpImageQuerySize) {
2895 /* The image coordinate is always 4 components but we may not have that
2896 * many. Swizzle to compensate.
2897 */
2898 intrin->src[1] = nir_src_for_ssa(expand_to_vec4(&b->nb, image.coord));
2899 intrin->src[2] = nir_src_for_ssa(image.sample);
2900 }
2901
2902 nir_intrinsic_set_access(intrin, image.image->access);
2903
2904 switch (opcode) {
2905 case SpvOpAtomicLoad:
2906 case SpvOpImageQuerySize:
2907 case SpvOpImageRead:
2908 if (opcode == SpvOpImageRead || opcode == SpvOpAtomicLoad) {
2909 /* Only OpImageRead can support a lod parameter if
2910 * SPV_AMD_shader_image_load_store_lod is used but the current NIR
2911 * intrinsics definition for atomics requires us to set it for
2912 * OpAtomicLoad.
2913 */
2914 intrin->src[3] = nir_src_for_ssa(image.lod);
2915 }
2916 break;
2917 case SpvOpAtomicStore:
2918 case SpvOpImageWrite: {
2919 const uint32_t value_id = opcode == SpvOpAtomicStore ? w[4] : w[3];
2920 nir_ssa_def *value = vtn_ssa_value(b, value_id)->def;
2921 /* nir_intrinsic_image_deref_store always takes a vec4 value */
2922 assert(op == nir_intrinsic_image_deref_store);
2923 intrin->num_components = 4;
2924 intrin->src[3] = nir_src_for_ssa(expand_to_vec4(&b->nb, value));
2925 /* Only OpImageWrite can support a lod parameter if
2926 * SPV_AMD_shader_image_load_store_lod is used but the current NIR
2927 * intrinsics definition for atomics requires us to set it for
2928 * OpAtomicStore.
2929 */
2930 intrin->src[4] = nir_src_for_ssa(image.lod);
2931 break;
2932 }
2933
2934 case SpvOpAtomicCompareExchange:
2935 case SpvOpAtomicCompareExchangeWeak:
2936 case SpvOpAtomicIIncrement:
2937 case SpvOpAtomicIDecrement:
2938 case SpvOpAtomicExchange:
2939 case SpvOpAtomicIAdd:
2940 case SpvOpAtomicISub:
2941 case SpvOpAtomicSMin:
2942 case SpvOpAtomicUMin:
2943 case SpvOpAtomicSMax:
2944 case SpvOpAtomicUMax:
2945 case SpvOpAtomicAnd:
2946 case SpvOpAtomicOr:
2947 case SpvOpAtomicXor:
2948 fill_common_atomic_sources(b, opcode, w, &intrin->src[3]);
2949 break;
2950
2951 default:
2952 vtn_fail_with_opcode("Invalid image opcode", opcode);
2953 }
2954
2955 /* Image operations implicitly have the Image storage memory semantics. */
2956 semantics |= SpvMemorySemanticsImageMemoryMask;
2957
2958 SpvMemorySemanticsMask before_semantics;
2959 SpvMemorySemanticsMask after_semantics;
2960 vtn_split_barrier_semantics(b, semantics, &before_semantics, &after_semantics);
2961
2962 if (before_semantics)
2963 vtn_emit_memory_barrier(b, scope, before_semantics);
2964
2965 if (opcode != SpvOpImageWrite && opcode != SpvOpAtomicStore) {
2966 struct vtn_type *type = vtn_value(b, w[1], vtn_value_type_type)->type;
2967
2968 unsigned dest_components = glsl_get_vector_elements(type->type);
2969 intrin->num_components = nir_intrinsic_infos[op].dest_components;
2970 if (intrin->num_components == 0)
2971 intrin->num_components = dest_components;
2972
2973 nir_ssa_dest_init(&intrin->instr, &intrin->dest,
2974 intrin->num_components, 32, NULL);
2975
2976 nir_builder_instr_insert(&b->nb, &intrin->instr);
2977
2978 nir_ssa_def *result = &intrin->dest.ssa;
2979 if (intrin->num_components != dest_components)
2980 result = nir_channels(&b->nb, result, (1 << dest_components) - 1);
2981
2982 struct vtn_value *val =
2983 vtn_push_ssa(b, w[2], type, vtn_create_ssa_value(b, type->type));
2984 val->ssa->def = result;
2985 } else {
2986 nir_builder_instr_insert(&b->nb, &intrin->instr);
2987 }
2988
2989 if (after_semantics)
2990 vtn_emit_memory_barrier(b, scope, after_semantics);
2991 }
2992
2993 static nir_intrinsic_op
2994 get_ssbo_nir_atomic_op(struct vtn_builder *b, SpvOp opcode)
2995 {
2996 switch (opcode) {
2997 case SpvOpAtomicLoad: return nir_intrinsic_load_ssbo;
2998 case SpvOpAtomicStore: return nir_intrinsic_store_ssbo;
2999 #define OP(S, N) case SpvOp##S: return nir_intrinsic_ssbo_##N;
3000 OP(AtomicExchange, atomic_exchange)
3001 OP(AtomicCompareExchange, atomic_comp_swap)
3002 OP(AtomicCompareExchangeWeak, atomic_comp_swap)
3003 OP(AtomicIIncrement, atomic_add)
3004 OP(AtomicIDecrement, atomic_add)
3005 OP(AtomicIAdd, atomic_add)
3006 OP(AtomicISub, atomic_add)
3007 OP(AtomicSMin, atomic_imin)
3008 OP(AtomicUMin, atomic_umin)
3009 OP(AtomicSMax, atomic_imax)
3010 OP(AtomicUMax, atomic_umax)
3011 OP(AtomicAnd, atomic_and)
3012 OP(AtomicOr, atomic_or)
3013 OP(AtomicXor, atomic_xor)
3014 #undef OP
3015 default:
3016 vtn_fail_with_opcode("Invalid SSBO atomic", opcode);
3017 }
3018 }
3019
3020 static nir_intrinsic_op
3021 get_uniform_nir_atomic_op(struct vtn_builder *b, SpvOp opcode)
3022 {
3023 switch (opcode) {
3024 #define OP(S, N) case SpvOp##S: return nir_intrinsic_atomic_counter_ ##N;
3025 OP(AtomicLoad, read_deref)
3026 OP(AtomicExchange, exchange)
3027 OP(AtomicCompareExchange, comp_swap)
3028 OP(AtomicCompareExchangeWeak, comp_swap)
3029 OP(AtomicIIncrement, inc_deref)
3030 OP(AtomicIDecrement, post_dec_deref)
3031 OP(AtomicIAdd, add_deref)
3032 OP(AtomicISub, add_deref)
3033 OP(AtomicUMin, min_deref)
3034 OP(AtomicUMax, max_deref)
3035 OP(AtomicAnd, and_deref)
3036 OP(AtomicOr, or_deref)
3037 OP(AtomicXor, xor_deref)
3038 #undef OP
3039 default:
3040 /* We left the following out: AtomicStore, AtomicSMin and
3041 * AtomicSmax. Right now there are not nir intrinsics for them. At this
3042 * moment Atomic Counter support is needed for ARB_spirv support, so is
3043 * only need to support GLSL Atomic Counters that are uints and don't
3044 * allow direct storage.
3045 */
3046 vtn_fail("Invalid uniform atomic");
3047 }
3048 }
3049
3050 static nir_intrinsic_op
3051 get_deref_nir_atomic_op(struct vtn_builder *b, SpvOp opcode)
3052 {
3053 switch (opcode) {
3054 case SpvOpAtomicLoad: return nir_intrinsic_load_deref;
3055 case SpvOpAtomicStore: return nir_intrinsic_store_deref;
3056 #define OP(S, N) case SpvOp##S: return nir_intrinsic_deref_##N;
3057 OP(AtomicExchange, atomic_exchange)
3058 OP(AtomicCompareExchange, atomic_comp_swap)
3059 OP(AtomicCompareExchangeWeak, atomic_comp_swap)
3060 OP(AtomicIIncrement, atomic_add)
3061 OP(AtomicIDecrement, atomic_add)
3062 OP(AtomicIAdd, atomic_add)
3063 OP(AtomicISub, atomic_add)
3064 OP(AtomicSMin, atomic_imin)
3065 OP(AtomicUMin, atomic_umin)
3066 OP(AtomicSMax, atomic_imax)
3067 OP(AtomicUMax, atomic_umax)
3068 OP(AtomicAnd, atomic_and)
3069 OP(AtomicOr, atomic_or)
3070 OP(AtomicXor, atomic_xor)
3071 #undef OP
3072 default:
3073 vtn_fail_with_opcode("Invalid shared atomic", opcode);
3074 }
3075 }
3076
3077 /*
3078 * Handles shared atomics, ssbo atomics and atomic counters.
3079 */
3080 static void
3081 vtn_handle_atomics(struct vtn_builder *b, SpvOp opcode,
3082 const uint32_t *w, UNUSED unsigned count)
3083 {
3084 struct vtn_pointer *ptr;
3085 nir_intrinsic_instr *atomic;
3086
3087 SpvScope scope = SpvScopeInvocation;
3088 SpvMemorySemanticsMask semantics = 0;
3089
3090 switch (opcode) {
3091 case SpvOpAtomicLoad:
3092 case SpvOpAtomicExchange:
3093 case SpvOpAtomicCompareExchange:
3094 case SpvOpAtomicCompareExchangeWeak:
3095 case SpvOpAtomicIIncrement:
3096 case SpvOpAtomicIDecrement:
3097 case SpvOpAtomicIAdd:
3098 case SpvOpAtomicISub:
3099 case SpvOpAtomicSMin:
3100 case SpvOpAtomicUMin:
3101 case SpvOpAtomicSMax:
3102 case SpvOpAtomicUMax:
3103 case SpvOpAtomicAnd:
3104 case SpvOpAtomicOr:
3105 case SpvOpAtomicXor:
3106 ptr = vtn_value(b, w[3], vtn_value_type_pointer)->pointer;
3107 scope = vtn_constant_uint(b, w[4]);
3108 semantics = vtn_constant_uint(b, w[5]);
3109 break;
3110
3111 case SpvOpAtomicStore:
3112 ptr = vtn_value(b, w[1], vtn_value_type_pointer)->pointer;
3113 scope = vtn_constant_uint(b, w[2]);
3114 semantics = vtn_constant_uint(b, w[3]);
3115 break;
3116
3117 default:
3118 vtn_fail_with_opcode("Invalid SPIR-V atomic", opcode);
3119 }
3120
3121 /* uniform as "atomic counter uniform" */
3122 if (ptr->mode == vtn_variable_mode_uniform) {
3123 nir_deref_instr *deref = vtn_pointer_to_deref(b, ptr);
3124 const struct glsl_type *deref_type = deref->type;
3125 nir_intrinsic_op op = get_uniform_nir_atomic_op(b, opcode);
3126 atomic = nir_intrinsic_instr_create(b->nb.shader, op);
3127 atomic->src[0] = nir_src_for_ssa(&deref->dest.ssa);
3128
3129 /* SSBO needs to initialize index/offset. In this case we don't need to,
3130 * as that info is already stored on the ptr->var->var nir_variable (see
3131 * vtn_create_variable)
3132 */
3133
3134 switch (opcode) {
3135 case SpvOpAtomicLoad:
3136 atomic->num_components = glsl_get_vector_elements(deref_type);
3137 break;
3138
3139 case SpvOpAtomicStore:
3140 atomic->num_components = glsl_get_vector_elements(deref_type);
3141 nir_intrinsic_set_write_mask(atomic, (1 << atomic->num_components) - 1);
3142 break;
3143
3144 case SpvOpAtomicExchange:
3145 case SpvOpAtomicCompareExchange:
3146 case SpvOpAtomicCompareExchangeWeak:
3147 case SpvOpAtomicIIncrement:
3148 case SpvOpAtomicIDecrement:
3149 case SpvOpAtomicIAdd:
3150 case SpvOpAtomicISub:
3151 case SpvOpAtomicSMin:
3152 case SpvOpAtomicUMin:
3153 case SpvOpAtomicSMax:
3154 case SpvOpAtomicUMax:
3155 case SpvOpAtomicAnd:
3156 case SpvOpAtomicOr:
3157 case SpvOpAtomicXor:
3158 /* Nothing: we don't need to call fill_common_atomic_sources here, as
3159 * atomic counter uniforms doesn't have sources
3160 */
3161 break;
3162
3163 default:
3164 unreachable("Invalid SPIR-V atomic");
3165
3166 }
3167 } else if (vtn_pointer_uses_ssa_offset(b, ptr)) {
3168 nir_ssa_def *offset, *index;
3169 offset = vtn_pointer_to_offset(b, ptr, &index);
3170
3171 assert(ptr->mode == vtn_variable_mode_ssbo);
3172
3173 nir_intrinsic_op op = get_ssbo_nir_atomic_op(b, opcode);
3174 atomic = nir_intrinsic_instr_create(b->nb.shader, op);
3175
3176 int src = 0;
3177 switch (opcode) {
3178 case SpvOpAtomicLoad:
3179 atomic->num_components = glsl_get_vector_elements(ptr->type->type);
3180 nir_intrinsic_set_align(atomic, 4, 0);
3181 if (ptr->mode == vtn_variable_mode_ssbo)
3182 atomic->src[src++] = nir_src_for_ssa(index);
3183 atomic->src[src++] = nir_src_for_ssa(offset);
3184 break;
3185
3186 case SpvOpAtomicStore:
3187 atomic->num_components = glsl_get_vector_elements(ptr->type->type);
3188 nir_intrinsic_set_write_mask(atomic, (1 << atomic->num_components) - 1);
3189 nir_intrinsic_set_align(atomic, 4, 0);
3190 atomic->src[src++] = nir_src_for_ssa(vtn_ssa_value(b, w[4])->def);
3191 if (ptr->mode == vtn_variable_mode_ssbo)
3192 atomic->src[src++] = nir_src_for_ssa(index);
3193 atomic->src[src++] = nir_src_for_ssa(offset);
3194 break;
3195
3196 case SpvOpAtomicExchange:
3197 case SpvOpAtomicCompareExchange:
3198 case SpvOpAtomicCompareExchangeWeak:
3199 case SpvOpAtomicIIncrement:
3200 case SpvOpAtomicIDecrement:
3201 case SpvOpAtomicIAdd:
3202 case SpvOpAtomicISub:
3203 case SpvOpAtomicSMin:
3204 case SpvOpAtomicUMin:
3205 case SpvOpAtomicSMax:
3206 case SpvOpAtomicUMax:
3207 case SpvOpAtomicAnd:
3208 case SpvOpAtomicOr:
3209 case SpvOpAtomicXor:
3210 if (ptr->mode == vtn_variable_mode_ssbo)
3211 atomic->src[src++] = nir_src_for_ssa(index);
3212 atomic->src[src++] = nir_src_for_ssa(offset);
3213 fill_common_atomic_sources(b, opcode, w, &atomic->src[src]);
3214 break;
3215
3216 default:
3217 vtn_fail_with_opcode("Invalid SPIR-V atomic", opcode);
3218 }
3219 } else {
3220 nir_deref_instr *deref = vtn_pointer_to_deref(b, ptr);
3221 const struct glsl_type *deref_type = deref->type;
3222 nir_intrinsic_op op = get_deref_nir_atomic_op(b, opcode);
3223 atomic = nir_intrinsic_instr_create(b->nb.shader, op);
3224 atomic->src[0] = nir_src_for_ssa(&deref->dest.ssa);
3225
3226 switch (opcode) {
3227 case SpvOpAtomicLoad:
3228 atomic->num_components = glsl_get_vector_elements(deref_type);
3229 break;
3230
3231 case SpvOpAtomicStore:
3232 atomic->num_components = glsl_get_vector_elements(deref_type);
3233 nir_intrinsic_set_write_mask(atomic, (1 << atomic->num_components) - 1);
3234 atomic->src[1] = nir_src_for_ssa(vtn_ssa_value(b, w[4])->def);
3235 break;
3236
3237 case SpvOpAtomicExchange:
3238 case SpvOpAtomicCompareExchange:
3239 case SpvOpAtomicCompareExchangeWeak:
3240 case SpvOpAtomicIIncrement:
3241 case SpvOpAtomicIDecrement:
3242 case SpvOpAtomicIAdd:
3243 case SpvOpAtomicISub:
3244 case SpvOpAtomicSMin:
3245 case SpvOpAtomicUMin:
3246 case SpvOpAtomicSMax:
3247 case SpvOpAtomicUMax:
3248 case SpvOpAtomicAnd:
3249 case SpvOpAtomicOr:
3250 case SpvOpAtomicXor:
3251 fill_common_atomic_sources(b, opcode, w, &atomic->src[1]);
3252 break;
3253
3254 default:
3255 vtn_fail_with_opcode("Invalid SPIR-V atomic", opcode);
3256 }
3257 }
3258
3259 /* Atomic ordering operations will implicitly apply to the atomic operation
3260 * storage class, so include that too.
3261 */
3262 semantics |= vtn_storage_class_to_memory_semantics(ptr->ptr_type->storage_class);
3263
3264 SpvMemorySemanticsMask before_semantics;
3265 SpvMemorySemanticsMask after_semantics;
3266 vtn_split_barrier_semantics(b, semantics, &before_semantics, &after_semantics);
3267
3268 if (before_semantics)
3269 vtn_emit_memory_barrier(b, scope, before_semantics);
3270
3271 if (opcode != SpvOpAtomicStore) {
3272 struct vtn_type *type = vtn_value(b, w[1], vtn_value_type_type)->type;
3273
3274 nir_ssa_dest_init(&atomic->instr, &atomic->dest,
3275 glsl_get_vector_elements(type->type),
3276 glsl_get_bit_size(type->type), NULL);
3277
3278 struct vtn_ssa_value *ssa = rzalloc(b, struct vtn_ssa_value);
3279 ssa->def = &atomic->dest.ssa;
3280 ssa->type = type->type;
3281 vtn_push_ssa(b, w[2], type, ssa);
3282 }
3283
3284 nir_builder_instr_insert(&b->nb, &atomic->instr);
3285
3286 if (after_semantics)
3287 vtn_emit_memory_barrier(b, scope, after_semantics);
3288 }
3289
3290 static nir_alu_instr *
3291 create_vec(struct vtn_builder *b, unsigned num_components, unsigned bit_size)
3292 {
3293 nir_op op = nir_op_vec(num_components);
3294 nir_alu_instr *vec = nir_alu_instr_create(b->shader, op);
3295 nir_ssa_dest_init(&vec->instr, &vec->dest.dest, num_components,
3296 bit_size, NULL);
3297 vec->dest.write_mask = (1 << num_components) - 1;
3298
3299 return vec;
3300 }
3301
3302 struct vtn_ssa_value *
3303 vtn_ssa_transpose(struct vtn_builder *b, struct vtn_ssa_value *src)
3304 {
3305 if (src->transposed)
3306 return src->transposed;
3307
3308 struct vtn_ssa_value *dest =
3309 vtn_create_ssa_value(b, glsl_transposed_type(src->type));
3310
3311 for (unsigned i = 0; i < glsl_get_matrix_columns(dest->type); i++) {
3312 nir_alu_instr *vec = create_vec(b, glsl_get_matrix_columns(src->type),
3313 glsl_get_bit_size(src->type));
3314 if (glsl_type_is_vector_or_scalar(src->type)) {
3315 vec->src[0].src = nir_src_for_ssa(src->def);
3316 vec->src[0].swizzle[0] = i;
3317 } else {
3318 for (unsigned j = 0; j < glsl_get_matrix_columns(src->type); j++) {
3319 vec->src[j].src = nir_src_for_ssa(src->elems[j]->def);
3320 vec->src[j].swizzle[0] = i;
3321 }
3322 }
3323 nir_builder_instr_insert(&b->nb, &vec->instr);
3324 dest->elems[i]->def = &vec->dest.dest.ssa;
3325 }
3326
3327 dest->transposed = src;
3328
3329 return dest;
3330 }
3331
3332 nir_ssa_def *
3333 vtn_vector_extract(struct vtn_builder *b, nir_ssa_def *src, unsigned index)
3334 {
3335 return nir_channel(&b->nb, src, index);
3336 }
3337
3338 nir_ssa_def *
3339 vtn_vector_insert(struct vtn_builder *b, nir_ssa_def *src, nir_ssa_def *insert,
3340 unsigned index)
3341 {
3342 nir_alu_instr *vec = create_vec(b, src->num_components,
3343 src->bit_size);
3344
3345 for (unsigned i = 0; i < src->num_components; i++) {
3346 if (i == index) {
3347 vec->src[i].src = nir_src_for_ssa(insert);
3348 } else {
3349 vec->src[i].src = nir_src_for_ssa(src);
3350 vec->src[i].swizzle[0] = i;
3351 }
3352 }
3353
3354 nir_builder_instr_insert(&b->nb, &vec->instr);
3355
3356 return &vec->dest.dest.ssa;
3357 }
3358
3359 static nir_ssa_def *
3360 nir_ieq_imm(nir_builder *b, nir_ssa_def *x, uint64_t i)
3361 {
3362 return nir_ieq(b, x, nir_imm_intN_t(b, i, x->bit_size));
3363 }
3364
3365 nir_ssa_def *
3366 vtn_vector_extract_dynamic(struct vtn_builder *b, nir_ssa_def *src,
3367 nir_ssa_def *index)
3368 {
3369 return nir_vector_extract(&b->nb, src, nir_i2i(&b->nb, index, 32));
3370 }
3371
3372 nir_ssa_def *
3373 vtn_vector_insert_dynamic(struct vtn_builder *b, nir_ssa_def *src,
3374 nir_ssa_def *insert, nir_ssa_def *index)
3375 {
3376 nir_ssa_def *dest = vtn_vector_insert(b, src, insert, 0);
3377 for (unsigned i = 1; i < src->num_components; i++)
3378 dest = nir_bcsel(&b->nb, nir_ieq_imm(&b->nb, index, i),
3379 vtn_vector_insert(b, src, insert, i), dest);
3380
3381 return dest;
3382 }
3383
3384 static nir_ssa_def *
3385 vtn_vector_shuffle(struct vtn_builder *b, unsigned num_components,
3386 nir_ssa_def *src0, nir_ssa_def *src1,
3387 const uint32_t *indices)
3388 {
3389 nir_alu_instr *vec = create_vec(b, num_components, src0->bit_size);
3390
3391 for (unsigned i = 0; i < num_components; i++) {
3392 uint32_t index = indices[i];
3393 if (index == 0xffffffff) {
3394 vec->src[i].src =
3395 nir_src_for_ssa(nir_ssa_undef(&b->nb, 1, src0->bit_size));
3396 } else if (index < src0->num_components) {
3397 vec->src[i].src = nir_src_for_ssa(src0);
3398 vec->src[i].swizzle[0] = index;
3399 } else {
3400 vec->src[i].src = nir_src_for_ssa(src1);
3401 vec->src[i].swizzle[0] = index - src0->num_components;
3402 }
3403 }
3404
3405 nir_builder_instr_insert(&b->nb, &vec->instr);
3406
3407 return &vec->dest.dest.ssa;
3408 }
3409
3410 /*
3411 * Concatentates a number of vectors/scalars together to produce a vector
3412 */
3413 static nir_ssa_def *
3414 vtn_vector_construct(struct vtn_builder *b, unsigned num_components,
3415 unsigned num_srcs, nir_ssa_def **srcs)
3416 {
3417 nir_alu_instr *vec = create_vec(b, num_components, srcs[0]->bit_size);
3418
3419 /* From the SPIR-V 1.1 spec for OpCompositeConstruct:
3420 *
3421 * "When constructing a vector, there must be at least two Constituent
3422 * operands."
3423 */
3424 vtn_assert(num_srcs >= 2);
3425
3426 unsigned dest_idx = 0;
3427 for (unsigned i = 0; i < num_srcs; i++) {
3428 nir_ssa_def *src = srcs[i];
3429 vtn_assert(dest_idx + src->num_components <= num_components);
3430 for (unsigned j = 0; j < src->num_components; j++) {
3431 vec->src[dest_idx].src = nir_src_for_ssa(src);
3432 vec->src[dest_idx].swizzle[0] = j;
3433 dest_idx++;
3434 }
3435 }
3436
3437 /* From the SPIR-V 1.1 spec for OpCompositeConstruct:
3438 *
3439 * "When constructing a vector, the total number of components in all
3440 * the operands must equal the number of components in Result Type."
3441 */
3442 vtn_assert(dest_idx == num_components);
3443
3444 nir_builder_instr_insert(&b->nb, &vec->instr);
3445
3446 return &vec->dest.dest.ssa;
3447 }
3448
3449 static struct vtn_ssa_value *
3450 vtn_composite_copy(void *mem_ctx, struct vtn_ssa_value *src)
3451 {
3452 struct vtn_ssa_value *dest = rzalloc(mem_ctx, struct vtn_ssa_value);
3453 dest->type = src->type;
3454
3455 if (glsl_type_is_vector_or_scalar(src->type)) {
3456 dest->def = src->def;
3457 } else {
3458 unsigned elems = glsl_get_length(src->type);
3459
3460 dest->elems = ralloc_array(mem_ctx, struct vtn_ssa_value *, elems);
3461 for (unsigned i = 0; i < elems; i++)
3462 dest->elems[i] = vtn_composite_copy(mem_ctx, src->elems[i]);
3463 }
3464
3465 return dest;
3466 }
3467
3468 static struct vtn_ssa_value *
3469 vtn_composite_insert(struct vtn_builder *b, struct vtn_ssa_value *src,
3470 struct vtn_ssa_value *insert, const uint32_t *indices,
3471 unsigned num_indices)
3472 {
3473 struct vtn_ssa_value *dest = vtn_composite_copy(b, src);
3474
3475 struct vtn_ssa_value *cur = dest;
3476 unsigned i;
3477 for (i = 0; i < num_indices - 1; i++) {
3478 cur = cur->elems[indices[i]];
3479 }
3480
3481 if (glsl_type_is_vector_or_scalar(cur->type)) {
3482 /* According to the SPIR-V spec, OpCompositeInsert may work down to
3483 * the component granularity. In that case, the last index will be
3484 * the index to insert the scalar into the vector.
3485 */
3486
3487 cur->def = vtn_vector_insert(b, cur->def, insert->def, indices[i]);
3488 } else {
3489 cur->elems[indices[i]] = insert;
3490 }
3491
3492 return dest;
3493 }
3494
3495 static struct vtn_ssa_value *
3496 vtn_composite_extract(struct vtn_builder *b, struct vtn_ssa_value *src,
3497 const uint32_t *indices, unsigned num_indices)
3498 {
3499 struct vtn_ssa_value *cur = src;
3500 for (unsigned i = 0; i < num_indices; i++) {
3501 if (glsl_type_is_vector_or_scalar(cur->type)) {
3502 vtn_assert(i == num_indices - 1);
3503 /* According to the SPIR-V spec, OpCompositeExtract may work down to
3504 * the component granularity. The last index will be the index of the
3505 * vector to extract.
3506 */
3507
3508 struct vtn_ssa_value *ret = rzalloc(b, struct vtn_ssa_value);
3509 ret->type = glsl_scalar_type(glsl_get_base_type(cur->type));
3510 ret->def = vtn_vector_extract(b, cur->def, indices[i]);
3511 return ret;
3512 } else {
3513 cur = cur->elems[indices[i]];
3514 }
3515 }
3516
3517 return cur;
3518 }
3519
3520 static void
3521 vtn_handle_composite(struct vtn_builder *b, SpvOp opcode,
3522 const uint32_t *w, unsigned count)
3523 {
3524 struct vtn_type *type = vtn_value(b, w[1], vtn_value_type_type)->type;
3525 struct vtn_ssa_value *ssa = vtn_create_ssa_value(b, type->type);
3526
3527 switch (opcode) {
3528 case SpvOpVectorExtractDynamic:
3529 ssa->def = vtn_vector_extract_dynamic(b, vtn_ssa_value(b, w[3])->def,
3530 vtn_ssa_value(b, w[4])->def);
3531 break;
3532
3533 case SpvOpVectorInsertDynamic:
3534 ssa->def = vtn_vector_insert_dynamic(b, vtn_ssa_value(b, w[3])->def,
3535 vtn_ssa_value(b, w[4])->def,
3536 vtn_ssa_value(b, w[5])->def);
3537 break;
3538
3539 case SpvOpVectorShuffle:
3540 ssa->def = vtn_vector_shuffle(b, glsl_get_vector_elements(type->type),
3541 vtn_ssa_value(b, w[3])->def,
3542 vtn_ssa_value(b, w[4])->def,
3543 w + 5);
3544 break;
3545
3546 case SpvOpCompositeConstruct: {
3547 unsigned elems = count - 3;
3548 assume(elems >= 1);
3549 if (glsl_type_is_vector_or_scalar(type->type)) {
3550 nir_ssa_def *srcs[NIR_MAX_VEC_COMPONENTS];
3551 for (unsigned i = 0; i < elems; i++)
3552 srcs[i] = vtn_ssa_value(b, w[3 + i])->def;
3553 ssa->def =
3554 vtn_vector_construct(b, glsl_get_vector_elements(type->type),
3555 elems, srcs);
3556 } else {
3557 ssa->elems = ralloc_array(b, struct vtn_ssa_value *, elems);
3558 for (unsigned i = 0; i < elems; i++)
3559 ssa->elems[i] = vtn_ssa_value(b, w[3 + i]);
3560 }
3561 break;
3562 }
3563 case SpvOpCompositeExtract:
3564 ssa = vtn_composite_extract(b, vtn_ssa_value(b, w[3]),
3565 w + 4, count - 4);
3566 break;
3567
3568 case SpvOpCompositeInsert:
3569 ssa = vtn_composite_insert(b, vtn_ssa_value(b, w[4]),
3570 vtn_ssa_value(b, w[3]),
3571 w + 5, count - 5);
3572 break;
3573
3574 case SpvOpCopyLogical:
3575 case SpvOpCopyObject:
3576 ssa = vtn_composite_copy(b, vtn_ssa_value(b, w[3]));
3577 break;
3578
3579 default:
3580 vtn_fail_with_opcode("unknown composite operation", opcode);
3581 }
3582
3583 vtn_push_ssa(b, w[2], type, ssa);
3584 }
3585
3586 static void
3587 vtn_emit_barrier(struct vtn_builder *b, nir_intrinsic_op op)
3588 {
3589 nir_intrinsic_instr *intrin = nir_intrinsic_instr_create(b->shader, op);
3590 nir_builder_instr_insert(&b->nb, &intrin->instr);
3591 }
3592
3593 void
3594 vtn_emit_memory_barrier(struct vtn_builder *b, SpvScope scope,
3595 SpvMemorySemanticsMask semantics)
3596 {
3597 if (b->options->use_scoped_memory_barrier) {
3598 vtn_emit_scoped_memory_barrier(b, scope, semantics);
3599 return;
3600 }
3601
3602 static const SpvMemorySemanticsMask all_memory_semantics =
3603 SpvMemorySemanticsUniformMemoryMask |
3604 SpvMemorySemanticsWorkgroupMemoryMask |
3605 SpvMemorySemanticsAtomicCounterMemoryMask |
3606 SpvMemorySemanticsImageMemoryMask;
3607
3608 /* If we're not actually doing a memory barrier, bail */
3609 if (!(semantics & all_memory_semantics))
3610 return;
3611
3612 /* GL and Vulkan don't have these */
3613 vtn_assert(scope != SpvScopeCrossDevice);
3614
3615 if (scope == SpvScopeSubgroup)
3616 return; /* Nothing to do here */
3617
3618 if (scope == SpvScopeWorkgroup) {
3619 vtn_emit_barrier(b, nir_intrinsic_group_memory_barrier);
3620 return;
3621 }
3622
3623 /* There's only two scopes thing left */
3624 vtn_assert(scope == SpvScopeInvocation || scope == SpvScopeDevice);
3625
3626 if ((semantics & all_memory_semantics) == all_memory_semantics) {
3627 vtn_emit_barrier(b, nir_intrinsic_memory_barrier);
3628 return;
3629 }
3630
3631 /* Issue a bunch of more specific barriers */
3632 uint32_t bits = semantics;
3633 while (bits) {
3634 SpvMemorySemanticsMask semantic = 1 << u_bit_scan(&bits);
3635 switch (semantic) {
3636 case SpvMemorySemanticsUniformMemoryMask:
3637 vtn_emit_barrier(b, nir_intrinsic_memory_barrier_buffer);
3638 break;
3639 case SpvMemorySemanticsWorkgroupMemoryMask:
3640 vtn_emit_barrier(b, nir_intrinsic_memory_barrier_shared);
3641 break;
3642 case SpvMemorySemanticsAtomicCounterMemoryMask:
3643 vtn_emit_barrier(b, nir_intrinsic_memory_barrier_atomic_counter);
3644 break;
3645 case SpvMemorySemanticsImageMemoryMask:
3646 vtn_emit_barrier(b, nir_intrinsic_memory_barrier_image);
3647 break;
3648 case SpvMemorySemanticsOutputMemoryMask:
3649 if (b->nb.shader->info.stage == MESA_SHADER_TESS_CTRL)
3650 vtn_emit_barrier(b, nir_intrinsic_memory_barrier_tcs_patch);
3651 break;
3652 default:
3653 break;;
3654 }
3655 }
3656 }
3657
3658 static void
3659 vtn_handle_barrier(struct vtn_builder *b, SpvOp opcode,
3660 const uint32_t *w, UNUSED unsigned count)
3661 {
3662 switch (opcode) {
3663 case SpvOpEmitVertex:
3664 case SpvOpEmitStreamVertex:
3665 case SpvOpEndPrimitive:
3666 case SpvOpEndStreamPrimitive: {
3667 nir_intrinsic_op intrinsic_op;
3668 switch (opcode) {
3669 case SpvOpEmitVertex:
3670 case SpvOpEmitStreamVertex:
3671 intrinsic_op = nir_intrinsic_emit_vertex;
3672 break;
3673 case SpvOpEndPrimitive:
3674 case SpvOpEndStreamPrimitive:
3675 intrinsic_op = nir_intrinsic_end_primitive;
3676 break;
3677 default:
3678 unreachable("Invalid opcode");
3679 }
3680
3681 nir_intrinsic_instr *intrin =
3682 nir_intrinsic_instr_create(b->shader, intrinsic_op);
3683
3684 switch (opcode) {
3685 case SpvOpEmitStreamVertex:
3686 case SpvOpEndStreamPrimitive: {
3687 unsigned stream = vtn_constant_uint(b, w[1]);
3688 nir_intrinsic_set_stream_id(intrin, stream);
3689 break;
3690 }
3691
3692 default:
3693 break;
3694 }
3695
3696 nir_builder_instr_insert(&b->nb, &intrin->instr);
3697 break;
3698 }
3699
3700 case SpvOpMemoryBarrier: {
3701 SpvScope scope = vtn_constant_uint(b, w[1]);
3702 SpvMemorySemanticsMask semantics = vtn_constant_uint(b, w[2]);
3703 vtn_emit_memory_barrier(b, scope, semantics);
3704 return;
3705 }
3706
3707 case SpvOpControlBarrier: {
3708 SpvScope execution_scope = vtn_constant_uint(b, w[1]);
3709 SpvScope memory_scope = vtn_constant_uint(b, w[2]);
3710 SpvMemorySemanticsMask memory_semantics = vtn_constant_uint(b, w[3]);
3711
3712 /* GLSLang, prior to commit 8297936dd6eb3, emitted OpControlBarrier with
3713 * memory semantics of None for GLSL barrier().
3714 */
3715 if (b->wa_glslang_cs_barrier &&
3716 b->nb.shader->info.stage == MESA_SHADER_COMPUTE &&
3717 execution_scope == SpvScopeWorkgroup &&
3718 memory_semantics == SpvMemorySemanticsMaskNone) {
3719 memory_scope = SpvScopeWorkgroup;
3720 memory_semantics = SpvMemorySemanticsAcquireReleaseMask |
3721 SpvMemorySemanticsWorkgroupMemoryMask;
3722 }
3723
3724 /* From the SPIR-V spec:
3725 *
3726 * "When used with the TessellationControl execution model, it also
3727 * implicitly synchronizes the Output Storage Class: Writes to Output
3728 * variables performed by any invocation executed prior to a
3729 * OpControlBarrier will be visible to any other invocation after
3730 * return from that OpControlBarrier."
3731 */
3732 if (b->nb.shader->info.stage == MESA_SHADER_TESS_CTRL) {
3733 memory_semantics &= ~(SpvMemorySemanticsAcquireMask |
3734 SpvMemorySemanticsReleaseMask |
3735 SpvMemorySemanticsAcquireReleaseMask |
3736 SpvMemorySemanticsSequentiallyConsistentMask);
3737 memory_semantics |= SpvMemorySemanticsAcquireReleaseMask |
3738 SpvMemorySemanticsOutputMemoryMask;
3739 }
3740
3741 vtn_emit_memory_barrier(b, memory_scope, memory_semantics);
3742
3743 if (execution_scope == SpvScopeWorkgroup)
3744 vtn_emit_barrier(b, nir_intrinsic_control_barrier);
3745 break;
3746 }
3747
3748 default:
3749 unreachable("unknown barrier instruction");
3750 }
3751 }
3752
3753 static unsigned
3754 gl_primitive_from_spv_execution_mode(struct vtn_builder *b,
3755 SpvExecutionMode mode)
3756 {
3757 switch (mode) {
3758 case SpvExecutionModeInputPoints:
3759 case SpvExecutionModeOutputPoints:
3760 return 0; /* GL_POINTS */
3761 case SpvExecutionModeInputLines:
3762 return 1; /* GL_LINES */
3763 case SpvExecutionModeInputLinesAdjacency:
3764 return 0x000A; /* GL_LINE_STRIP_ADJACENCY_ARB */
3765 case SpvExecutionModeTriangles:
3766 return 4; /* GL_TRIANGLES */
3767 case SpvExecutionModeInputTrianglesAdjacency:
3768 return 0x000C; /* GL_TRIANGLES_ADJACENCY_ARB */
3769 case SpvExecutionModeQuads:
3770 return 7; /* GL_QUADS */
3771 case SpvExecutionModeIsolines:
3772 return 0x8E7A; /* GL_ISOLINES */
3773 case SpvExecutionModeOutputLineStrip:
3774 return 3; /* GL_LINE_STRIP */
3775 case SpvExecutionModeOutputTriangleStrip:
3776 return 5; /* GL_TRIANGLE_STRIP */
3777 default:
3778 vtn_fail("Invalid primitive type: %s (%u)",
3779 spirv_executionmode_to_string(mode), mode);
3780 }
3781 }
3782
3783 static unsigned
3784 vertices_in_from_spv_execution_mode(struct vtn_builder *b,
3785 SpvExecutionMode mode)
3786 {
3787 switch (mode) {
3788 case SpvExecutionModeInputPoints:
3789 return 1;
3790 case SpvExecutionModeInputLines:
3791 return 2;
3792 case SpvExecutionModeInputLinesAdjacency:
3793 return 4;
3794 case SpvExecutionModeTriangles:
3795 return 3;
3796 case SpvExecutionModeInputTrianglesAdjacency:
3797 return 6;
3798 default:
3799 vtn_fail("Invalid GS input mode: %s (%u)",
3800 spirv_executionmode_to_string(mode), mode);
3801 }
3802 }
3803
3804 static gl_shader_stage
3805 stage_for_execution_model(struct vtn_builder *b, SpvExecutionModel model)
3806 {
3807 switch (model) {
3808 case SpvExecutionModelVertex:
3809 return MESA_SHADER_VERTEX;
3810 case SpvExecutionModelTessellationControl:
3811 return MESA_SHADER_TESS_CTRL;
3812 case SpvExecutionModelTessellationEvaluation:
3813 return MESA_SHADER_TESS_EVAL;
3814 case SpvExecutionModelGeometry:
3815 return MESA_SHADER_GEOMETRY;
3816 case SpvExecutionModelFragment:
3817 return MESA_SHADER_FRAGMENT;
3818 case SpvExecutionModelGLCompute:
3819 return MESA_SHADER_COMPUTE;
3820 case SpvExecutionModelKernel:
3821 return MESA_SHADER_KERNEL;
3822 default:
3823 vtn_fail("Unsupported execution model: %s (%u)",
3824 spirv_executionmodel_to_string(model), model);
3825 }
3826 }
3827
3828 #define spv_check_supported(name, cap) do { \
3829 if (!(b->options && b->options->caps.name)) \
3830 vtn_warn("Unsupported SPIR-V capability: %s (%u)", \
3831 spirv_capability_to_string(cap), cap); \
3832 } while(0)
3833
3834
3835 void
3836 vtn_handle_entry_point(struct vtn_builder *b, const uint32_t *w,
3837 unsigned count)
3838 {
3839 struct vtn_value *entry_point = &b->values[w[2]];
3840 /* Let this be a name label regardless */
3841 unsigned name_words;
3842 entry_point->name = vtn_string_literal(b, &w[3], count - 3, &name_words);
3843
3844 if (strcmp(entry_point->name, b->entry_point_name) != 0 ||
3845 stage_for_execution_model(b, w[1]) != b->entry_point_stage)
3846 return;
3847
3848 vtn_assert(b->entry_point == NULL);
3849 b->entry_point = entry_point;
3850 }
3851
3852 static bool
3853 vtn_handle_preamble_instruction(struct vtn_builder *b, SpvOp opcode,
3854 const uint32_t *w, unsigned count)
3855 {
3856 switch (opcode) {
3857 case SpvOpSource: {
3858 const char *lang;
3859 switch (w[1]) {
3860 default:
3861 case SpvSourceLanguageUnknown: lang = "unknown"; break;
3862 case SpvSourceLanguageESSL: lang = "ESSL"; break;
3863 case SpvSourceLanguageGLSL: lang = "GLSL"; break;
3864 case SpvSourceLanguageOpenCL_C: lang = "OpenCL C"; break;
3865 case SpvSourceLanguageOpenCL_CPP: lang = "OpenCL C++"; break;
3866 case SpvSourceLanguageHLSL: lang = "HLSL"; break;
3867 }
3868
3869 uint32_t version = w[2];
3870
3871 const char *file =
3872 (count > 3) ? vtn_value(b, w[3], vtn_value_type_string)->str : "";
3873
3874 vtn_info("Parsing SPIR-V from %s %u source file %s", lang, version, file);
3875 break;
3876 }
3877
3878 case SpvOpSourceExtension:
3879 case SpvOpSourceContinued:
3880 case SpvOpExtension:
3881 case SpvOpModuleProcessed:
3882 /* Unhandled, but these are for debug so that's ok. */
3883 break;
3884
3885 case SpvOpCapability: {
3886 SpvCapability cap = w[1];
3887 switch (cap) {
3888 case SpvCapabilityMatrix:
3889 case SpvCapabilityShader:
3890 case SpvCapabilityGeometry:
3891 case SpvCapabilityGeometryPointSize:
3892 case SpvCapabilityUniformBufferArrayDynamicIndexing:
3893 case SpvCapabilitySampledImageArrayDynamicIndexing:
3894 case SpvCapabilityStorageBufferArrayDynamicIndexing:
3895 case SpvCapabilityStorageImageArrayDynamicIndexing:
3896 case SpvCapabilityImageRect:
3897 case SpvCapabilitySampledRect:
3898 case SpvCapabilitySampled1D:
3899 case SpvCapabilityImage1D:
3900 case SpvCapabilitySampledCubeArray:
3901 case SpvCapabilityImageCubeArray:
3902 case SpvCapabilitySampledBuffer:
3903 case SpvCapabilityImageBuffer:
3904 case SpvCapabilityImageQuery:
3905 case SpvCapabilityDerivativeControl:
3906 case SpvCapabilityInterpolationFunction:
3907 case SpvCapabilityMultiViewport:
3908 case SpvCapabilitySampleRateShading:
3909 case SpvCapabilityClipDistance:
3910 case SpvCapabilityCullDistance:
3911 case SpvCapabilityInputAttachment:
3912 case SpvCapabilityImageGatherExtended:
3913 case SpvCapabilityStorageImageExtendedFormats:
3914 case SpvCapabilityVector16:
3915 break;
3916
3917 case SpvCapabilityLinkage:
3918 case SpvCapabilityFloat16Buffer:
3919 case SpvCapabilitySparseResidency:
3920 vtn_warn("Unsupported SPIR-V capability: %s",
3921 spirv_capability_to_string(cap));
3922 break;
3923
3924 case SpvCapabilityMinLod:
3925 spv_check_supported(min_lod, cap);
3926 break;
3927
3928 case SpvCapabilityAtomicStorage:
3929 spv_check_supported(atomic_storage, cap);
3930 break;
3931
3932 case SpvCapabilityFloat64:
3933 spv_check_supported(float64, cap);
3934 break;
3935 case SpvCapabilityInt64:
3936 spv_check_supported(int64, cap);
3937 break;
3938 case SpvCapabilityInt16:
3939 spv_check_supported(int16, cap);
3940 break;
3941 case SpvCapabilityInt8:
3942 spv_check_supported(int8, cap);
3943 break;
3944
3945 case SpvCapabilityTransformFeedback:
3946 spv_check_supported(transform_feedback, cap);
3947 break;
3948
3949 case SpvCapabilityGeometryStreams:
3950 spv_check_supported(geometry_streams, cap);
3951 break;
3952
3953 case SpvCapabilityInt64Atomics:
3954 spv_check_supported(int64_atomics, cap);
3955 break;
3956
3957 case SpvCapabilityStorageImageMultisample:
3958 spv_check_supported(storage_image_ms, cap);
3959 break;
3960
3961 case SpvCapabilityAddresses:
3962 spv_check_supported(address, cap);
3963 break;
3964
3965 case SpvCapabilityKernel:
3966 spv_check_supported(kernel, cap);
3967 break;
3968
3969 case SpvCapabilityImageBasic:
3970 case SpvCapabilityImageReadWrite:
3971 case SpvCapabilityImageMipmap:
3972 case SpvCapabilityPipes:
3973 case SpvCapabilityDeviceEnqueue:
3974 case SpvCapabilityLiteralSampler:
3975 case SpvCapabilityGenericPointer:
3976 vtn_warn("Unsupported OpenCL-style SPIR-V capability: %s",
3977 spirv_capability_to_string(cap));
3978 break;
3979
3980 case SpvCapabilityImageMSArray:
3981 spv_check_supported(image_ms_array, cap);
3982 break;
3983
3984 case SpvCapabilityTessellation:
3985 case SpvCapabilityTessellationPointSize:
3986 spv_check_supported(tessellation, cap);
3987 break;
3988
3989 case SpvCapabilityDrawParameters:
3990 spv_check_supported(draw_parameters, cap);
3991 break;
3992
3993 case SpvCapabilityStorageImageReadWithoutFormat:
3994 spv_check_supported(image_read_without_format, cap);
3995 break;
3996
3997 case SpvCapabilityStorageImageWriteWithoutFormat:
3998 spv_check_supported(image_write_without_format, cap);
3999 break;
4000
4001 case SpvCapabilityDeviceGroup:
4002 spv_check_supported(device_group, cap);
4003 break;
4004
4005 case SpvCapabilityMultiView:
4006 spv_check_supported(multiview, cap);
4007 break;
4008
4009 case SpvCapabilityGroupNonUniform:
4010 spv_check_supported(subgroup_basic, cap);
4011 break;
4012
4013 case SpvCapabilitySubgroupVoteKHR:
4014 case SpvCapabilityGroupNonUniformVote:
4015 spv_check_supported(subgroup_vote, cap);
4016 break;
4017
4018 case SpvCapabilitySubgroupBallotKHR:
4019 case SpvCapabilityGroupNonUniformBallot:
4020 spv_check_supported(subgroup_ballot, cap);
4021 break;
4022
4023 case SpvCapabilityGroupNonUniformShuffle:
4024 case SpvCapabilityGroupNonUniformShuffleRelative:
4025 spv_check_supported(subgroup_shuffle, cap);
4026 break;
4027
4028 case SpvCapabilityGroupNonUniformQuad:
4029 spv_check_supported(subgroup_quad, cap);
4030 break;
4031
4032 case SpvCapabilityGroupNonUniformArithmetic:
4033 case SpvCapabilityGroupNonUniformClustered:
4034 spv_check_supported(subgroup_arithmetic, cap);
4035 break;
4036
4037 case SpvCapabilityGroups:
4038 spv_check_supported(amd_shader_ballot, cap);
4039 break;
4040
4041 case SpvCapabilityVariablePointersStorageBuffer:
4042 case SpvCapabilityVariablePointers:
4043 spv_check_supported(variable_pointers, cap);
4044 b->variable_pointers = true;
4045 break;
4046
4047 case SpvCapabilityStorageUniformBufferBlock16:
4048 case SpvCapabilityStorageUniform16:
4049 case SpvCapabilityStoragePushConstant16:
4050 case SpvCapabilityStorageInputOutput16:
4051 spv_check_supported(storage_16bit, cap);
4052 break;
4053
4054 case SpvCapabilityShaderLayer:
4055 case SpvCapabilityShaderViewportIndex:
4056 case SpvCapabilityShaderViewportIndexLayerEXT:
4057 spv_check_supported(shader_viewport_index_layer, cap);
4058 break;
4059
4060 case SpvCapabilityStorageBuffer8BitAccess:
4061 case SpvCapabilityUniformAndStorageBuffer8BitAccess:
4062 case SpvCapabilityStoragePushConstant8:
4063 spv_check_supported(storage_8bit, cap);
4064 break;
4065
4066 case SpvCapabilityShaderNonUniformEXT:
4067 spv_check_supported(descriptor_indexing, cap);
4068 break;
4069
4070 case SpvCapabilityInputAttachmentArrayDynamicIndexingEXT:
4071 case SpvCapabilityUniformTexelBufferArrayDynamicIndexingEXT:
4072 case SpvCapabilityStorageTexelBufferArrayDynamicIndexingEXT:
4073 spv_check_supported(descriptor_array_dynamic_indexing, cap);
4074 break;
4075
4076 case SpvCapabilityUniformBufferArrayNonUniformIndexingEXT:
4077 case SpvCapabilitySampledImageArrayNonUniformIndexingEXT:
4078 case SpvCapabilityStorageBufferArrayNonUniformIndexingEXT:
4079 case SpvCapabilityStorageImageArrayNonUniformIndexingEXT:
4080 case SpvCapabilityInputAttachmentArrayNonUniformIndexingEXT:
4081 case SpvCapabilityUniformTexelBufferArrayNonUniformIndexingEXT:
4082 case SpvCapabilityStorageTexelBufferArrayNonUniformIndexingEXT:
4083 spv_check_supported(descriptor_array_non_uniform_indexing, cap);
4084 break;
4085
4086 case SpvCapabilityRuntimeDescriptorArrayEXT:
4087 spv_check_supported(runtime_descriptor_array, cap);
4088 break;
4089
4090 case SpvCapabilityStencilExportEXT:
4091 spv_check_supported(stencil_export, cap);
4092 break;
4093
4094 case SpvCapabilitySampleMaskPostDepthCoverage:
4095 spv_check_supported(post_depth_coverage, cap);
4096 break;
4097
4098 case SpvCapabilityDenormFlushToZero:
4099 case SpvCapabilityDenormPreserve:
4100 case SpvCapabilitySignedZeroInfNanPreserve:
4101 case SpvCapabilityRoundingModeRTE:
4102 case SpvCapabilityRoundingModeRTZ:
4103 spv_check_supported(float_controls, cap);
4104 break;
4105
4106 case SpvCapabilityPhysicalStorageBufferAddresses:
4107 spv_check_supported(physical_storage_buffer_address, cap);
4108 break;
4109
4110 case SpvCapabilityComputeDerivativeGroupQuadsNV:
4111 case SpvCapabilityComputeDerivativeGroupLinearNV:
4112 spv_check_supported(derivative_group, cap);
4113 break;
4114
4115 case SpvCapabilityFloat16:
4116 spv_check_supported(float16, cap);
4117 break;
4118
4119 case SpvCapabilityFragmentShaderSampleInterlockEXT:
4120 spv_check_supported(fragment_shader_sample_interlock, cap);
4121 break;
4122
4123 case SpvCapabilityFragmentShaderPixelInterlockEXT:
4124 spv_check_supported(fragment_shader_pixel_interlock, cap);
4125 break;
4126
4127 case SpvCapabilityDemoteToHelperInvocationEXT:
4128 spv_check_supported(demote_to_helper_invocation, cap);
4129 break;
4130
4131 case SpvCapabilityShaderClockKHR:
4132 spv_check_supported(shader_clock, cap);
4133 break;
4134
4135 case SpvCapabilityVulkanMemoryModel:
4136 spv_check_supported(vk_memory_model, cap);
4137 break;
4138
4139 case SpvCapabilityVulkanMemoryModelDeviceScope:
4140 spv_check_supported(vk_memory_model_device_scope, cap);
4141 break;
4142
4143 case SpvCapabilityImageReadWriteLodAMD:
4144 spv_check_supported(amd_image_read_write_lod, cap);
4145 break;
4146
4147 case SpvCapabilityIntegerFunctions2INTEL:
4148 spv_check_supported(integer_functions2, cap);
4149 break;
4150
4151 case SpvCapabilityFragmentMaskAMD:
4152 spv_check_supported(amd_fragment_mask, cap);
4153 break;
4154
4155 default:
4156 vtn_fail("Unhandled capability: %s (%u)",
4157 spirv_capability_to_string(cap), cap);
4158 }
4159 break;
4160 }
4161
4162 case SpvOpExtInstImport:
4163 vtn_handle_extension(b, opcode, w, count);
4164 break;
4165
4166 case SpvOpMemoryModel:
4167 switch (w[1]) {
4168 case SpvAddressingModelPhysical32:
4169 vtn_fail_if(b->shader->info.stage != MESA_SHADER_KERNEL,
4170 "AddressingModelPhysical32 only supported for kernels");
4171 b->shader->info.cs.ptr_size = 32;
4172 b->physical_ptrs = true;
4173 b->options->shared_addr_format = nir_address_format_32bit_global;
4174 b->options->global_addr_format = nir_address_format_32bit_global;
4175 b->options->temp_addr_format = nir_address_format_32bit_global;
4176 break;
4177 case SpvAddressingModelPhysical64:
4178 vtn_fail_if(b->shader->info.stage != MESA_SHADER_KERNEL,
4179 "AddressingModelPhysical64 only supported for kernels");
4180 b->shader->info.cs.ptr_size = 64;
4181 b->physical_ptrs = true;
4182 b->options->shared_addr_format = nir_address_format_64bit_global;
4183 b->options->global_addr_format = nir_address_format_64bit_global;
4184 b->options->temp_addr_format = nir_address_format_64bit_global;
4185 break;
4186 case SpvAddressingModelLogical:
4187 vtn_fail_if(b->shader->info.stage >= MESA_SHADER_STAGES,
4188 "AddressingModelLogical only supported for shaders");
4189 b->physical_ptrs = false;
4190 break;
4191 case SpvAddressingModelPhysicalStorageBuffer64:
4192 vtn_fail_if(!b->options ||
4193 !b->options->caps.physical_storage_buffer_address,
4194 "AddressingModelPhysicalStorageBuffer64 not supported");
4195 break;
4196 default:
4197 vtn_fail("Unknown addressing model: %s (%u)",
4198 spirv_addressingmodel_to_string(w[1]), w[1]);
4199 break;
4200 }
4201
4202 switch (w[2]) {
4203 case SpvMemoryModelSimple:
4204 case SpvMemoryModelGLSL450:
4205 case SpvMemoryModelOpenCL:
4206 break;
4207 case SpvMemoryModelVulkan:
4208 vtn_fail_if(!b->options->caps.vk_memory_model,
4209 "Vulkan memory model is unsupported by this driver");
4210 break;
4211 default:
4212 vtn_fail("Unsupported memory model: %s",
4213 spirv_memorymodel_to_string(w[2]));
4214 break;
4215 }
4216 break;
4217
4218 case SpvOpEntryPoint:
4219 vtn_handle_entry_point(b, w, count);
4220 break;
4221
4222 case SpvOpString:
4223 vtn_push_value(b, w[1], vtn_value_type_string)->str =
4224 vtn_string_literal(b, &w[2], count - 2, NULL);
4225 break;
4226
4227 case SpvOpName:
4228 b->values[w[1]].name = vtn_string_literal(b, &w[2], count - 2, NULL);
4229 break;
4230
4231 case SpvOpMemberName:
4232 /* TODO */
4233 break;
4234
4235 case SpvOpExecutionMode:
4236 case SpvOpExecutionModeId:
4237 case SpvOpDecorationGroup:
4238 case SpvOpDecorate:
4239 case SpvOpDecorateId:
4240 case SpvOpMemberDecorate:
4241 case SpvOpGroupDecorate:
4242 case SpvOpGroupMemberDecorate:
4243 case SpvOpDecorateString:
4244 case SpvOpMemberDecorateString:
4245 vtn_handle_decoration(b, opcode, w, count);
4246 break;
4247
4248 case SpvOpExtInst: {
4249 struct vtn_value *val = vtn_value(b, w[3], vtn_value_type_extension);
4250 if (val->ext_handler == vtn_handle_non_semantic_instruction) {
4251 /* NonSemantic extended instructions are acceptable in preamble. */
4252 vtn_handle_non_semantic_instruction(b, w[4], w, count);
4253 return true;
4254 } else {
4255 return false; /* End of preamble. */
4256 }
4257 }
4258
4259 default:
4260 return false; /* End of preamble */
4261 }
4262
4263 return true;
4264 }
4265
4266 static void
4267 vtn_handle_execution_mode(struct vtn_builder *b, struct vtn_value *entry_point,
4268 const struct vtn_decoration *mode, UNUSED void *data)
4269 {
4270 vtn_assert(b->entry_point == entry_point);
4271
4272 switch(mode->exec_mode) {
4273 case SpvExecutionModeOriginUpperLeft:
4274 case SpvExecutionModeOriginLowerLeft:
4275 vtn_assert(b->shader->info.stage == MESA_SHADER_FRAGMENT);
4276 b->shader->info.fs.origin_upper_left =
4277 (mode->exec_mode == SpvExecutionModeOriginUpperLeft);
4278 break;
4279
4280 case SpvExecutionModeEarlyFragmentTests:
4281 vtn_assert(b->shader->info.stage == MESA_SHADER_FRAGMENT);
4282 b->shader->info.fs.early_fragment_tests = true;
4283 break;
4284
4285 case SpvExecutionModePostDepthCoverage:
4286 vtn_assert(b->shader->info.stage == MESA_SHADER_FRAGMENT);
4287 b->shader->info.fs.post_depth_coverage = true;
4288 break;
4289
4290 case SpvExecutionModeInvocations:
4291 vtn_assert(b->shader->info.stage == MESA_SHADER_GEOMETRY);
4292 b->shader->info.gs.invocations = MAX2(1, mode->operands[0]);
4293 break;
4294
4295 case SpvExecutionModeDepthReplacing:
4296 vtn_assert(b->shader->info.stage == MESA_SHADER_FRAGMENT);
4297 b->shader->info.fs.depth_layout = FRAG_DEPTH_LAYOUT_ANY;
4298 break;
4299 case SpvExecutionModeDepthGreater:
4300 vtn_assert(b->shader->info.stage == MESA_SHADER_FRAGMENT);
4301 b->shader->info.fs.depth_layout = FRAG_DEPTH_LAYOUT_GREATER;
4302 break;
4303 case SpvExecutionModeDepthLess:
4304 vtn_assert(b->shader->info.stage == MESA_SHADER_FRAGMENT);
4305 b->shader->info.fs.depth_layout = FRAG_DEPTH_LAYOUT_LESS;
4306 break;
4307 case SpvExecutionModeDepthUnchanged:
4308 vtn_assert(b->shader->info.stage == MESA_SHADER_FRAGMENT);
4309 b->shader->info.fs.depth_layout = FRAG_DEPTH_LAYOUT_UNCHANGED;
4310 break;
4311
4312 case SpvExecutionModeLocalSize:
4313 vtn_assert(gl_shader_stage_is_compute(b->shader->info.stage));
4314 b->shader->info.cs.local_size[0] = mode->operands[0];
4315 b->shader->info.cs.local_size[1] = mode->operands[1];
4316 b->shader->info.cs.local_size[2] = mode->operands[2];
4317 break;
4318
4319 case SpvExecutionModeLocalSizeId:
4320 b->shader->info.cs.local_size[0] = vtn_constant_uint(b, mode->operands[0]);
4321 b->shader->info.cs.local_size[1] = vtn_constant_uint(b, mode->operands[1]);
4322 b->shader->info.cs.local_size[2] = vtn_constant_uint(b, mode->operands[2]);
4323 break;
4324
4325 case SpvExecutionModeLocalSizeHint:
4326 case SpvExecutionModeLocalSizeHintId:
4327 break; /* Nothing to do with this */
4328
4329 case SpvExecutionModeOutputVertices:
4330 if (b->shader->info.stage == MESA_SHADER_TESS_CTRL ||
4331 b->shader->info.stage == MESA_SHADER_TESS_EVAL) {
4332 b->shader->info.tess.tcs_vertices_out = mode->operands[0];
4333 } else {
4334 vtn_assert(b->shader->info.stage == MESA_SHADER_GEOMETRY);
4335 b->shader->info.gs.vertices_out = mode->operands[0];
4336 }
4337 break;
4338
4339 case SpvExecutionModeInputPoints:
4340 case SpvExecutionModeInputLines:
4341 case SpvExecutionModeInputLinesAdjacency:
4342 case SpvExecutionModeTriangles:
4343 case SpvExecutionModeInputTrianglesAdjacency:
4344 case SpvExecutionModeQuads:
4345 case SpvExecutionModeIsolines:
4346 if (b->shader->info.stage == MESA_SHADER_TESS_CTRL ||
4347 b->shader->info.stage == MESA_SHADER_TESS_EVAL) {
4348 b->shader->info.tess.primitive_mode =
4349 gl_primitive_from_spv_execution_mode(b, mode->exec_mode);
4350 } else {
4351 vtn_assert(b->shader->info.stage == MESA_SHADER_GEOMETRY);
4352 b->shader->info.gs.vertices_in =
4353 vertices_in_from_spv_execution_mode(b, mode->exec_mode);
4354 b->shader->info.gs.input_primitive =
4355 gl_primitive_from_spv_execution_mode(b, mode->exec_mode);
4356 }
4357 break;
4358
4359 case SpvExecutionModeOutputPoints:
4360 case SpvExecutionModeOutputLineStrip:
4361 case SpvExecutionModeOutputTriangleStrip:
4362 vtn_assert(b->shader->info.stage == MESA_SHADER_GEOMETRY);
4363 b->shader->info.gs.output_primitive =
4364 gl_primitive_from_spv_execution_mode(b, mode->exec_mode);
4365 break;
4366
4367 case SpvExecutionModeSpacingEqual:
4368 vtn_assert(b->shader->info.stage == MESA_SHADER_TESS_CTRL ||
4369 b->shader->info.stage == MESA_SHADER_TESS_EVAL);
4370 b->shader->info.tess.spacing = TESS_SPACING_EQUAL;
4371 break;
4372 case SpvExecutionModeSpacingFractionalEven:
4373 vtn_assert(b->shader->info.stage == MESA_SHADER_TESS_CTRL ||
4374 b->shader->info.stage == MESA_SHADER_TESS_EVAL);
4375 b->shader->info.tess.spacing = TESS_SPACING_FRACTIONAL_EVEN;
4376 break;
4377 case SpvExecutionModeSpacingFractionalOdd:
4378 vtn_assert(b->shader->info.stage == MESA_SHADER_TESS_CTRL ||
4379 b->shader->info.stage == MESA_SHADER_TESS_EVAL);
4380 b->shader->info.tess.spacing = TESS_SPACING_FRACTIONAL_ODD;
4381 break;
4382 case SpvExecutionModeVertexOrderCw:
4383 vtn_assert(b->shader->info.stage == MESA_SHADER_TESS_CTRL ||
4384 b->shader->info.stage == MESA_SHADER_TESS_EVAL);
4385 b->shader->info.tess.ccw = false;
4386 break;
4387 case SpvExecutionModeVertexOrderCcw:
4388 vtn_assert(b->shader->info.stage == MESA_SHADER_TESS_CTRL ||
4389 b->shader->info.stage == MESA_SHADER_TESS_EVAL);
4390 b->shader->info.tess.ccw = true;
4391 break;
4392 case SpvExecutionModePointMode:
4393 vtn_assert(b->shader->info.stage == MESA_SHADER_TESS_CTRL ||
4394 b->shader->info.stage == MESA_SHADER_TESS_EVAL);
4395 b->shader->info.tess.point_mode = true;
4396 break;
4397
4398 case SpvExecutionModePixelCenterInteger:
4399 vtn_assert(b->shader->info.stage == MESA_SHADER_FRAGMENT);
4400 b->shader->info.fs.pixel_center_integer = true;
4401 break;
4402
4403 case SpvExecutionModeXfb:
4404 b->shader->info.has_transform_feedback_varyings = true;
4405 break;
4406
4407 case SpvExecutionModeVecTypeHint:
4408 break; /* OpenCL */
4409
4410 case SpvExecutionModeContractionOff:
4411 if (b->shader->info.stage != MESA_SHADER_KERNEL)
4412 vtn_warn("ExectionMode only allowed for CL-style kernels: %s",
4413 spirv_executionmode_to_string(mode->exec_mode));
4414 else
4415 b->exact = true;
4416 break;
4417
4418 case SpvExecutionModeStencilRefReplacingEXT:
4419 vtn_assert(b->shader->info.stage == MESA_SHADER_FRAGMENT);
4420 break;
4421
4422 case SpvExecutionModeDerivativeGroupQuadsNV:
4423 vtn_assert(b->shader->info.stage == MESA_SHADER_COMPUTE);
4424 b->shader->info.cs.derivative_group = DERIVATIVE_GROUP_QUADS;
4425 break;
4426
4427 case SpvExecutionModeDerivativeGroupLinearNV:
4428 vtn_assert(b->shader->info.stage == MESA_SHADER_COMPUTE);
4429 b->shader->info.cs.derivative_group = DERIVATIVE_GROUP_LINEAR;
4430 break;
4431
4432 case SpvExecutionModePixelInterlockOrderedEXT:
4433 vtn_assert(b->shader->info.stage == MESA_SHADER_FRAGMENT);
4434 b->shader->info.fs.pixel_interlock_ordered = true;
4435 break;
4436
4437 case SpvExecutionModePixelInterlockUnorderedEXT:
4438 vtn_assert(b->shader->info.stage == MESA_SHADER_FRAGMENT);
4439 b->shader->info.fs.pixel_interlock_unordered = true;
4440 break;
4441
4442 case SpvExecutionModeSampleInterlockOrderedEXT:
4443 vtn_assert(b->shader->info.stage == MESA_SHADER_FRAGMENT);
4444 b->shader->info.fs.sample_interlock_ordered = true;
4445 break;
4446
4447 case SpvExecutionModeSampleInterlockUnorderedEXT:
4448 vtn_assert(b->shader->info.stage == MESA_SHADER_FRAGMENT);
4449 b->shader->info.fs.sample_interlock_unordered = true;
4450 break;
4451
4452 case SpvExecutionModeDenormPreserve:
4453 case SpvExecutionModeDenormFlushToZero:
4454 case SpvExecutionModeSignedZeroInfNanPreserve:
4455 case SpvExecutionModeRoundingModeRTE:
4456 case SpvExecutionModeRoundingModeRTZ:
4457 /* Already handled in vtn_handle_rounding_mode_in_execution_mode() */
4458 break;
4459
4460 default:
4461 vtn_fail("Unhandled execution mode: %s (%u)",
4462 spirv_executionmode_to_string(mode->exec_mode),
4463 mode->exec_mode);
4464 }
4465 }
4466
4467 static void
4468 vtn_handle_rounding_mode_in_execution_mode(struct vtn_builder *b, struct vtn_value *entry_point,
4469 const struct vtn_decoration *mode, void *data)
4470 {
4471 vtn_assert(b->entry_point == entry_point);
4472
4473 unsigned execution_mode = 0;
4474
4475 switch(mode->exec_mode) {
4476 case SpvExecutionModeDenormPreserve:
4477 switch (mode->operands[0]) {
4478 case 16: execution_mode = FLOAT_CONTROLS_DENORM_PRESERVE_FP16; break;
4479 case 32: execution_mode = FLOAT_CONTROLS_DENORM_PRESERVE_FP32; break;
4480 case 64: execution_mode = FLOAT_CONTROLS_DENORM_PRESERVE_FP64; break;
4481 default: vtn_fail("Floating point type not supported");
4482 }
4483 break;
4484 case SpvExecutionModeDenormFlushToZero:
4485 switch (mode->operands[0]) {
4486 case 16: execution_mode = FLOAT_CONTROLS_DENORM_FLUSH_TO_ZERO_FP16; break;
4487 case 32: execution_mode = FLOAT_CONTROLS_DENORM_FLUSH_TO_ZERO_FP32; break;
4488 case 64: execution_mode = FLOAT_CONTROLS_DENORM_FLUSH_TO_ZERO_FP64; break;
4489 default: vtn_fail("Floating point type not supported");
4490 }
4491 break;
4492 case SpvExecutionModeSignedZeroInfNanPreserve:
4493 switch (mode->operands[0]) {
4494 case 16: execution_mode = FLOAT_CONTROLS_SIGNED_ZERO_INF_NAN_PRESERVE_FP16; break;
4495 case 32: execution_mode = FLOAT_CONTROLS_SIGNED_ZERO_INF_NAN_PRESERVE_FP32; break;
4496 case 64: execution_mode = FLOAT_CONTROLS_SIGNED_ZERO_INF_NAN_PRESERVE_FP64; break;
4497 default: vtn_fail("Floating point type not supported");
4498 }
4499 break;
4500 case SpvExecutionModeRoundingModeRTE:
4501 switch (mode->operands[0]) {
4502 case 16: execution_mode = FLOAT_CONTROLS_ROUNDING_MODE_RTE_FP16; break;
4503 case 32: execution_mode = FLOAT_CONTROLS_ROUNDING_MODE_RTE_FP32; break;
4504 case 64: execution_mode = FLOAT_CONTROLS_ROUNDING_MODE_RTE_FP64; break;
4505 default: vtn_fail("Floating point type not supported");
4506 }
4507 break;
4508 case SpvExecutionModeRoundingModeRTZ:
4509 switch (mode->operands[0]) {
4510 case 16: execution_mode = FLOAT_CONTROLS_ROUNDING_MODE_RTZ_FP16; break;
4511 case 32: execution_mode = FLOAT_CONTROLS_ROUNDING_MODE_RTZ_FP32; break;
4512 case 64: execution_mode = FLOAT_CONTROLS_ROUNDING_MODE_RTZ_FP64; break;
4513 default: vtn_fail("Floating point type not supported");
4514 }
4515 break;
4516
4517 default:
4518 break;
4519 }
4520
4521 b->shader->info.float_controls_execution_mode |= execution_mode;
4522 }
4523
4524 static bool
4525 vtn_handle_variable_or_type_instruction(struct vtn_builder *b, SpvOp opcode,
4526 const uint32_t *w, unsigned count)
4527 {
4528 vtn_set_instruction_result_type(b, opcode, w, count);
4529
4530 switch (opcode) {
4531 case SpvOpSource:
4532 case SpvOpSourceContinued:
4533 case SpvOpSourceExtension:
4534 case SpvOpExtension:
4535 case SpvOpCapability:
4536 case SpvOpExtInstImport:
4537 case SpvOpMemoryModel:
4538 case SpvOpEntryPoint:
4539 case SpvOpExecutionMode:
4540 case SpvOpString:
4541 case SpvOpName:
4542 case SpvOpMemberName:
4543 case SpvOpDecorationGroup:
4544 case SpvOpDecorate:
4545 case SpvOpDecorateId:
4546 case SpvOpMemberDecorate:
4547 case SpvOpGroupDecorate:
4548 case SpvOpGroupMemberDecorate:
4549 case SpvOpDecorateString:
4550 case SpvOpMemberDecorateString:
4551 vtn_fail("Invalid opcode types and variables section");
4552 break;
4553
4554 case SpvOpTypeVoid:
4555 case SpvOpTypeBool:
4556 case SpvOpTypeInt:
4557 case SpvOpTypeFloat:
4558 case SpvOpTypeVector:
4559 case SpvOpTypeMatrix:
4560 case SpvOpTypeImage:
4561 case SpvOpTypeSampler:
4562 case SpvOpTypeSampledImage:
4563 case SpvOpTypeArray:
4564 case SpvOpTypeRuntimeArray:
4565 case SpvOpTypeStruct:
4566 case SpvOpTypeOpaque:
4567 case SpvOpTypePointer:
4568 case SpvOpTypeForwardPointer:
4569 case SpvOpTypeFunction:
4570 case SpvOpTypeEvent:
4571 case SpvOpTypeDeviceEvent:
4572 case SpvOpTypeReserveId:
4573 case SpvOpTypeQueue:
4574 case SpvOpTypePipe:
4575 vtn_handle_type(b, opcode, w, count);
4576 break;
4577
4578 case SpvOpConstantTrue:
4579 case SpvOpConstantFalse:
4580 case SpvOpConstant:
4581 case SpvOpConstantComposite:
4582 case SpvOpConstantSampler:
4583 case SpvOpConstantNull:
4584 case SpvOpSpecConstantTrue:
4585 case SpvOpSpecConstantFalse:
4586 case SpvOpSpecConstant:
4587 case SpvOpSpecConstantComposite:
4588 case SpvOpSpecConstantOp:
4589 vtn_handle_constant(b, opcode, w, count);
4590 break;
4591
4592 case SpvOpUndef:
4593 case SpvOpVariable:
4594 vtn_handle_variables(b, opcode, w, count);
4595 break;
4596
4597 case SpvOpExtInst: {
4598 struct vtn_value *val = vtn_value(b, w[3], vtn_value_type_extension);
4599 /* NonSemantic extended instructions are acceptable in preamble, others
4600 * will indicate the end of preamble.
4601 */
4602 return val->ext_handler == vtn_handle_non_semantic_instruction;
4603 }
4604
4605 default:
4606 return false; /* End of preamble */
4607 }
4608
4609 return true;
4610 }
4611
4612 static struct vtn_ssa_value *
4613 vtn_nir_select(struct vtn_builder *b, struct vtn_ssa_value *src0,
4614 struct vtn_ssa_value *src1, struct vtn_ssa_value *src2)
4615 {
4616 struct vtn_ssa_value *dest = rzalloc(b, struct vtn_ssa_value);
4617 dest->type = src1->type;
4618
4619 if (glsl_type_is_vector_or_scalar(src1->type)) {
4620 dest->def = nir_bcsel(&b->nb, src0->def, src1->def, src2->def);
4621 } else {
4622 unsigned elems = glsl_get_length(src1->type);
4623
4624 dest->elems = ralloc_array(b, struct vtn_ssa_value *, elems);
4625 for (unsigned i = 0; i < elems; i++) {
4626 dest->elems[i] = vtn_nir_select(b, src0,
4627 src1->elems[i], src2->elems[i]);
4628 }
4629 }
4630
4631 return dest;
4632 }
4633
4634 static void
4635 vtn_handle_select(struct vtn_builder *b, SpvOp opcode,
4636 const uint32_t *w, unsigned count)
4637 {
4638 /* Handle OpSelect up-front here because it needs to be able to handle
4639 * pointers and not just regular vectors and scalars.
4640 */
4641 struct vtn_value *res_val = vtn_untyped_value(b, w[2]);
4642 struct vtn_value *cond_val = vtn_untyped_value(b, w[3]);
4643 struct vtn_value *obj1_val = vtn_untyped_value(b, w[4]);
4644 struct vtn_value *obj2_val = vtn_untyped_value(b, w[5]);
4645
4646 vtn_fail_if(obj1_val->type != res_val->type ||
4647 obj2_val->type != res_val->type,
4648 "Object types must match the result type in OpSelect");
4649
4650 vtn_fail_if((cond_val->type->base_type != vtn_base_type_scalar &&
4651 cond_val->type->base_type != vtn_base_type_vector) ||
4652 !glsl_type_is_boolean(cond_val->type->type),
4653 "OpSelect must have either a vector of booleans or "
4654 "a boolean as Condition type");
4655
4656 vtn_fail_if(cond_val->type->base_type == vtn_base_type_vector &&
4657 (res_val->type->base_type != vtn_base_type_vector ||
4658 res_val->type->length != cond_val->type->length),
4659 "When Condition type in OpSelect is a vector, the Result "
4660 "type must be a vector of the same length");
4661
4662 switch (res_val->type->base_type) {
4663 case vtn_base_type_scalar:
4664 case vtn_base_type_vector:
4665 case vtn_base_type_matrix:
4666 case vtn_base_type_array:
4667 case vtn_base_type_struct:
4668 /* OK. */
4669 break;
4670 case vtn_base_type_pointer:
4671 /* We need to have actual storage for pointer types. */
4672 vtn_fail_if(res_val->type->type == NULL,
4673 "Invalid pointer result type for OpSelect");
4674 break;
4675 default:
4676 vtn_fail("Result type of OpSelect must be a scalar, composite, or pointer");
4677 }
4678
4679 struct vtn_type *res_type = vtn_value(b, w[1], vtn_value_type_type)->type;
4680 struct vtn_ssa_value *ssa = vtn_nir_select(b,
4681 vtn_ssa_value(b, w[3]), vtn_ssa_value(b, w[4]), vtn_ssa_value(b, w[5]));
4682
4683 vtn_push_ssa(b, w[2], res_type, ssa);
4684 }
4685
4686 static void
4687 vtn_handle_ptr(struct vtn_builder *b, SpvOp opcode,
4688 const uint32_t *w, unsigned count)
4689 {
4690 struct vtn_type *type1 = vtn_untyped_value(b, w[3])->type;
4691 struct vtn_type *type2 = vtn_untyped_value(b, w[4])->type;
4692 vtn_fail_if(type1->base_type != vtn_base_type_pointer ||
4693 type2->base_type != vtn_base_type_pointer,
4694 "%s operands must have pointer types",
4695 spirv_op_to_string(opcode));
4696 vtn_fail_if(type1->storage_class != type2->storage_class,
4697 "%s operands must have the same storage class",
4698 spirv_op_to_string(opcode));
4699
4700 struct vtn_type *vtn_type =
4701 vtn_value(b, w[1], vtn_value_type_type)->type;
4702 const struct glsl_type *type = vtn_type->type;
4703
4704 nir_address_format addr_format = vtn_mode_to_address_format(
4705 b, vtn_storage_class_to_mode(b, type1->storage_class, NULL, NULL));
4706
4707 nir_ssa_def *def;
4708
4709 switch (opcode) {
4710 case SpvOpPtrDiff: {
4711 /* OpPtrDiff returns the difference in number of elements (not byte offset). */
4712 unsigned elem_size, elem_align;
4713 glsl_get_natural_size_align_bytes(type1->deref->type,
4714 &elem_size, &elem_align);
4715
4716 def = nir_build_addr_isub(&b->nb,
4717 vtn_ssa_value(b, w[3])->def,
4718 vtn_ssa_value(b, w[4])->def,
4719 addr_format);
4720 def = nir_idiv(&b->nb, def, nir_imm_intN_t(&b->nb, elem_size, def->bit_size));
4721 def = nir_i2i(&b->nb, def, glsl_get_bit_size(type));
4722 break;
4723 }
4724
4725 case SpvOpPtrEqual:
4726 case SpvOpPtrNotEqual: {
4727 def = nir_build_addr_ieq(&b->nb,
4728 vtn_ssa_value(b, w[3])->def,
4729 vtn_ssa_value(b, w[4])->def,
4730 addr_format);
4731 if (opcode == SpvOpPtrNotEqual)
4732 def = nir_inot(&b->nb, def);
4733 break;
4734 }
4735
4736 default:
4737 unreachable("Invalid ptr operation");
4738 }
4739
4740 struct vtn_ssa_value *ssa_value = vtn_create_ssa_value(b, type);
4741 ssa_value->def = def;
4742 vtn_push_ssa(b, w[2], vtn_type, ssa_value);
4743 }
4744
4745 static bool
4746 vtn_handle_body_instruction(struct vtn_builder *b, SpvOp opcode,
4747 const uint32_t *w, unsigned count)
4748 {
4749 switch (opcode) {
4750 case SpvOpLabel:
4751 break;
4752
4753 case SpvOpLoopMerge:
4754 case SpvOpSelectionMerge:
4755 /* This is handled by cfg pre-pass and walk_blocks */
4756 break;
4757
4758 case SpvOpUndef: {
4759 struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_undef);
4760 val->type = vtn_value(b, w[1], vtn_value_type_type)->type;
4761 break;
4762 }
4763
4764 case SpvOpExtInst:
4765 vtn_handle_extension(b, opcode, w, count);
4766 break;
4767
4768 case SpvOpVariable:
4769 case SpvOpLoad:
4770 case SpvOpStore:
4771 case SpvOpCopyMemory:
4772 case SpvOpCopyMemorySized:
4773 case SpvOpAccessChain:
4774 case SpvOpPtrAccessChain:
4775 case SpvOpInBoundsAccessChain:
4776 case SpvOpInBoundsPtrAccessChain:
4777 case SpvOpArrayLength:
4778 case SpvOpConvertPtrToU:
4779 case SpvOpConvertUToPtr:
4780 vtn_handle_variables(b, opcode, w, count);
4781 break;
4782
4783 case SpvOpFunctionCall:
4784 vtn_handle_function_call(b, opcode, w, count);
4785 break;
4786
4787 case SpvOpSampledImage:
4788 case SpvOpImage:
4789 case SpvOpImageSampleImplicitLod:
4790 case SpvOpImageSampleExplicitLod:
4791 case SpvOpImageSampleDrefImplicitLod:
4792 case SpvOpImageSampleDrefExplicitLod:
4793 case SpvOpImageSampleProjImplicitLod:
4794 case SpvOpImageSampleProjExplicitLod:
4795 case SpvOpImageSampleProjDrefImplicitLod:
4796 case SpvOpImageSampleProjDrefExplicitLod:
4797 case SpvOpImageFetch:
4798 case SpvOpImageGather:
4799 case SpvOpImageDrefGather:
4800 case SpvOpImageQuerySizeLod:
4801 case SpvOpImageQueryLod:
4802 case SpvOpImageQueryLevels:
4803 case SpvOpImageQuerySamples:
4804 vtn_handle_texture(b, opcode, w, count);
4805 break;
4806
4807 case SpvOpImageRead:
4808 case SpvOpImageWrite:
4809 case SpvOpImageTexelPointer:
4810 vtn_handle_image(b, opcode, w, count);
4811 break;
4812
4813 case SpvOpImageQuerySize: {
4814 struct vtn_pointer *image =
4815 vtn_value(b, w[3], vtn_value_type_pointer)->pointer;
4816 if (glsl_type_is_image(image->type->type)) {
4817 vtn_handle_image(b, opcode, w, count);
4818 } else {
4819 vtn_assert(glsl_type_is_sampler(image->type->type));
4820 vtn_handle_texture(b, opcode, w, count);
4821 }
4822 break;
4823 }
4824
4825 case SpvOpFragmentMaskFetchAMD:
4826 case SpvOpFragmentFetchAMD:
4827 vtn_handle_texture(b, opcode, w, count);
4828 break;
4829
4830 case SpvOpAtomicLoad:
4831 case SpvOpAtomicExchange:
4832 case SpvOpAtomicCompareExchange:
4833 case SpvOpAtomicCompareExchangeWeak:
4834 case SpvOpAtomicIIncrement:
4835 case SpvOpAtomicIDecrement:
4836 case SpvOpAtomicIAdd:
4837 case SpvOpAtomicISub:
4838 case SpvOpAtomicSMin:
4839 case SpvOpAtomicUMin:
4840 case SpvOpAtomicSMax:
4841 case SpvOpAtomicUMax:
4842 case SpvOpAtomicAnd:
4843 case SpvOpAtomicOr:
4844 case SpvOpAtomicXor: {
4845 struct vtn_value *pointer = vtn_untyped_value(b, w[3]);
4846 if (pointer->value_type == vtn_value_type_image_pointer) {
4847 vtn_handle_image(b, opcode, w, count);
4848 } else {
4849 vtn_assert(pointer->value_type == vtn_value_type_pointer);
4850 vtn_handle_atomics(b, opcode, w, count);
4851 }
4852 break;
4853 }
4854
4855 case SpvOpAtomicStore: {
4856 struct vtn_value *pointer = vtn_untyped_value(b, w[1]);
4857 if (pointer->value_type == vtn_value_type_image_pointer) {
4858 vtn_handle_image(b, opcode, w, count);
4859 } else {
4860 vtn_assert(pointer->value_type == vtn_value_type_pointer);
4861 vtn_handle_atomics(b, opcode, w, count);
4862 }
4863 break;
4864 }
4865
4866 case SpvOpSelect:
4867 vtn_handle_select(b, opcode, w, count);
4868 break;
4869
4870 case SpvOpSNegate:
4871 case SpvOpFNegate:
4872 case SpvOpNot:
4873 case SpvOpAny:
4874 case SpvOpAll:
4875 case SpvOpConvertFToU:
4876 case SpvOpConvertFToS:
4877 case SpvOpConvertSToF:
4878 case SpvOpConvertUToF:
4879 case SpvOpUConvert:
4880 case SpvOpSConvert:
4881 case SpvOpFConvert:
4882 case SpvOpQuantizeToF16:
4883 case SpvOpPtrCastToGeneric:
4884 case SpvOpGenericCastToPtr:
4885 case SpvOpIsNan:
4886 case SpvOpIsInf:
4887 case SpvOpIsFinite:
4888 case SpvOpIsNormal:
4889 case SpvOpSignBitSet:
4890 case SpvOpLessOrGreater:
4891 case SpvOpOrdered:
4892 case SpvOpUnordered:
4893 case SpvOpIAdd:
4894 case SpvOpFAdd:
4895 case SpvOpISub:
4896 case SpvOpFSub:
4897 case SpvOpIMul:
4898 case SpvOpFMul:
4899 case SpvOpUDiv:
4900 case SpvOpSDiv:
4901 case SpvOpFDiv:
4902 case SpvOpUMod:
4903 case SpvOpSRem:
4904 case SpvOpSMod:
4905 case SpvOpFRem:
4906 case SpvOpFMod:
4907 case SpvOpVectorTimesScalar:
4908 case SpvOpDot:
4909 case SpvOpIAddCarry:
4910 case SpvOpISubBorrow:
4911 case SpvOpUMulExtended:
4912 case SpvOpSMulExtended:
4913 case SpvOpShiftRightLogical:
4914 case SpvOpShiftRightArithmetic:
4915 case SpvOpShiftLeftLogical:
4916 case SpvOpLogicalEqual:
4917 case SpvOpLogicalNotEqual:
4918 case SpvOpLogicalOr:
4919 case SpvOpLogicalAnd:
4920 case SpvOpLogicalNot:
4921 case SpvOpBitwiseOr:
4922 case SpvOpBitwiseXor:
4923 case SpvOpBitwiseAnd:
4924 case SpvOpIEqual:
4925 case SpvOpFOrdEqual:
4926 case SpvOpFUnordEqual:
4927 case SpvOpINotEqual:
4928 case SpvOpFOrdNotEqual:
4929 case SpvOpFUnordNotEqual:
4930 case SpvOpULessThan:
4931 case SpvOpSLessThan:
4932 case SpvOpFOrdLessThan:
4933 case SpvOpFUnordLessThan:
4934 case SpvOpUGreaterThan:
4935 case SpvOpSGreaterThan:
4936 case SpvOpFOrdGreaterThan:
4937 case SpvOpFUnordGreaterThan:
4938 case SpvOpULessThanEqual:
4939 case SpvOpSLessThanEqual:
4940 case SpvOpFOrdLessThanEqual:
4941 case SpvOpFUnordLessThanEqual:
4942 case SpvOpUGreaterThanEqual:
4943 case SpvOpSGreaterThanEqual:
4944 case SpvOpFOrdGreaterThanEqual:
4945 case SpvOpFUnordGreaterThanEqual:
4946 case SpvOpDPdx:
4947 case SpvOpDPdy:
4948 case SpvOpFwidth:
4949 case SpvOpDPdxFine:
4950 case SpvOpDPdyFine:
4951 case SpvOpFwidthFine:
4952 case SpvOpDPdxCoarse:
4953 case SpvOpDPdyCoarse:
4954 case SpvOpFwidthCoarse:
4955 case SpvOpBitFieldInsert:
4956 case SpvOpBitFieldSExtract:
4957 case SpvOpBitFieldUExtract:
4958 case SpvOpBitReverse:
4959 case SpvOpBitCount:
4960 case SpvOpTranspose:
4961 case SpvOpOuterProduct:
4962 case SpvOpMatrixTimesScalar:
4963 case SpvOpVectorTimesMatrix:
4964 case SpvOpMatrixTimesVector:
4965 case SpvOpMatrixTimesMatrix:
4966 case SpvOpUCountLeadingZerosINTEL:
4967 case SpvOpUCountTrailingZerosINTEL:
4968 case SpvOpAbsISubINTEL:
4969 case SpvOpAbsUSubINTEL:
4970 case SpvOpIAddSatINTEL:
4971 case SpvOpUAddSatINTEL:
4972 case SpvOpIAverageINTEL:
4973 case SpvOpUAverageINTEL:
4974 case SpvOpIAverageRoundedINTEL:
4975 case SpvOpUAverageRoundedINTEL:
4976 case SpvOpISubSatINTEL:
4977 case SpvOpUSubSatINTEL:
4978 case SpvOpIMul32x16INTEL:
4979 case SpvOpUMul32x16INTEL:
4980 vtn_handle_alu(b, opcode, w, count);
4981 break;
4982
4983 case SpvOpBitcast:
4984 vtn_handle_bitcast(b, w, count);
4985 break;
4986
4987 case SpvOpVectorExtractDynamic:
4988 case SpvOpVectorInsertDynamic:
4989 case SpvOpVectorShuffle:
4990 case SpvOpCompositeConstruct:
4991 case SpvOpCompositeExtract:
4992 case SpvOpCompositeInsert:
4993 case SpvOpCopyLogical:
4994 case SpvOpCopyObject:
4995 vtn_handle_composite(b, opcode, w, count);
4996 break;
4997
4998 case SpvOpEmitVertex:
4999 case SpvOpEndPrimitive:
5000 case SpvOpEmitStreamVertex:
5001 case SpvOpEndStreamPrimitive:
5002 case SpvOpControlBarrier:
5003 case SpvOpMemoryBarrier:
5004 vtn_handle_barrier(b, opcode, w, count);
5005 break;
5006
5007 case SpvOpGroupNonUniformElect:
5008 case SpvOpGroupNonUniformAll:
5009 case SpvOpGroupNonUniformAny:
5010 case SpvOpGroupNonUniformAllEqual:
5011 case SpvOpGroupNonUniformBroadcast:
5012 case SpvOpGroupNonUniformBroadcastFirst:
5013 case SpvOpGroupNonUniformBallot:
5014 case SpvOpGroupNonUniformInverseBallot:
5015 case SpvOpGroupNonUniformBallotBitExtract:
5016 case SpvOpGroupNonUniformBallotBitCount:
5017 case SpvOpGroupNonUniformBallotFindLSB:
5018 case SpvOpGroupNonUniformBallotFindMSB:
5019 case SpvOpGroupNonUniformShuffle:
5020 case SpvOpGroupNonUniformShuffleXor:
5021 case SpvOpGroupNonUniformShuffleUp:
5022 case SpvOpGroupNonUniformShuffleDown:
5023 case SpvOpGroupNonUniformIAdd:
5024 case SpvOpGroupNonUniformFAdd:
5025 case SpvOpGroupNonUniformIMul:
5026 case SpvOpGroupNonUniformFMul:
5027 case SpvOpGroupNonUniformSMin:
5028 case SpvOpGroupNonUniformUMin:
5029 case SpvOpGroupNonUniformFMin:
5030 case SpvOpGroupNonUniformSMax:
5031 case SpvOpGroupNonUniformUMax:
5032 case SpvOpGroupNonUniformFMax:
5033 case SpvOpGroupNonUniformBitwiseAnd:
5034 case SpvOpGroupNonUniformBitwiseOr:
5035 case SpvOpGroupNonUniformBitwiseXor:
5036 case SpvOpGroupNonUniformLogicalAnd:
5037 case SpvOpGroupNonUniformLogicalOr:
5038 case SpvOpGroupNonUniformLogicalXor:
5039 case SpvOpGroupNonUniformQuadBroadcast:
5040 case SpvOpGroupNonUniformQuadSwap:
5041 case SpvOpGroupAll:
5042 case SpvOpGroupAny:
5043 case SpvOpGroupBroadcast:
5044 case SpvOpGroupIAdd:
5045 case SpvOpGroupFAdd:
5046 case SpvOpGroupFMin:
5047 case SpvOpGroupUMin:
5048 case SpvOpGroupSMin:
5049 case SpvOpGroupFMax:
5050 case SpvOpGroupUMax:
5051 case SpvOpGroupSMax:
5052 case SpvOpSubgroupBallotKHR:
5053 case SpvOpSubgroupFirstInvocationKHR:
5054 case SpvOpSubgroupReadInvocationKHR:
5055 case SpvOpSubgroupAllKHR:
5056 case SpvOpSubgroupAnyKHR:
5057 case SpvOpSubgroupAllEqualKHR:
5058 case SpvOpGroupIAddNonUniformAMD:
5059 case SpvOpGroupFAddNonUniformAMD:
5060 case SpvOpGroupFMinNonUniformAMD:
5061 case SpvOpGroupUMinNonUniformAMD:
5062 case SpvOpGroupSMinNonUniformAMD:
5063 case SpvOpGroupFMaxNonUniformAMD:
5064 case SpvOpGroupUMaxNonUniformAMD:
5065 case SpvOpGroupSMaxNonUniformAMD:
5066 vtn_handle_subgroup(b, opcode, w, count);
5067 break;
5068
5069 case SpvOpPtrDiff:
5070 case SpvOpPtrEqual:
5071 case SpvOpPtrNotEqual:
5072 vtn_handle_ptr(b, opcode, w, count);
5073 break;
5074
5075 case SpvOpBeginInvocationInterlockEXT:
5076 vtn_emit_barrier(b, nir_intrinsic_begin_invocation_interlock);
5077 break;
5078
5079 case SpvOpEndInvocationInterlockEXT:
5080 vtn_emit_barrier(b, nir_intrinsic_end_invocation_interlock);
5081 break;
5082
5083 case SpvOpDemoteToHelperInvocationEXT: {
5084 nir_intrinsic_instr *intrin =
5085 nir_intrinsic_instr_create(b->shader, nir_intrinsic_demote);
5086 nir_builder_instr_insert(&b->nb, &intrin->instr);
5087 break;
5088 }
5089
5090 case SpvOpIsHelperInvocationEXT: {
5091 nir_intrinsic_instr *intrin =
5092 nir_intrinsic_instr_create(b->shader, nir_intrinsic_is_helper_invocation);
5093 nir_ssa_dest_init(&intrin->instr, &intrin->dest, 1, 1, NULL);
5094 nir_builder_instr_insert(&b->nb, &intrin->instr);
5095
5096 struct vtn_type *res_type =
5097 vtn_value(b, w[1], vtn_value_type_type)->type;
5098 struct vtn_ssa_value *val = vtn_create_ssa_value(b, res_type->type);
5099 val->def = &intrin->dest.ssa;
5100
5101 vtn_push_ssa(b, w[2], res_type, val);
5102 break;
5103 }
5104
5105 case SpvOpReadClockKHR: {
5106 assert(vtn_constant_uint(b, w[3]) == SpvScopeSubgroup);
5107
5108 /* Operation supports two result types: uvec2 and uint64_t. The NIR
5109 * intrinsic gives uvec2, so pack the result for the other case.
5110 */
5111 nir_intrinsic_instr *intrin =
5112 nir_intrinsic_instr_create(b->nb.shader, nir_intrinsic_shader_clock);
5113 nir_ssa_dest_init(&intrin->instr, &intrin->dest, 2, 32, NULL);
5114 nir_builder_instr_insert(&b->nb, &intrin->instr);
5115
5116 struct vtn_type *type = vtn_value(b, w[1], vtn_value_type_type)->type;
5117 const struct glsl_type *dest_type = type->type;
5118 nir_ssa_def *result;
5119
5120 if (glsl_type_is_vector(dest_type)) {
5121 assert(dest_type == glsl_vector_type(GLSL_TYPE_UINT, 2));
5122 result = &intrin->dest.ssa;
5123 } else {
5124 assert(glsl_type_is_scalar(dest_type));
5125 assert(glsl_get_base_type(dest_type) == GLSL_TYPE_UINT64);
5126 result = nir_pack_64_2x32(&b->nb, &intrin->dest.ssa);
5127 }
5128
5129 struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_ssa);
5130 val->type = type;
5131 val->ssa = vtn_create_ssa_value(b, dest_type);
5132 val->ssa->def = result;
5133 break;
5134 }
5135
5136 case SpvOpLifetimeStart:
5137 case SpvOpLifetimeStop:
5138 break;
5139
5140 default:
5141 vtn_fail_with_opcode("Unhandled opcode", opcode);
5142 }
5143
5144 return true;
5145 }
5146
5147 struct vtn_builder*
5148 vtn_create_builder(const uint32_t *words, size_t word_count,
5149 gl_shader_stage stage, const char *entry_point_name,
5150 const struct spirv_to_nir_options *options)
5151 {
5152 /* Initialize the vtn_builder object */
5153 struct vtn_builder *b = rzalloc(NULL, struct vtn_builder);
5154 struct spirv_to_nir_options *dup_options =
5155 ralloc(b, struct spirv_to_nir_options);
5156 *dup_options = *options;
5157
5158 b->spirv = words;
5159 b->spirv_word_count = word_count;
5160 b->file = NULL;
5161 b->line = -1;
5162 b->col = -1;
5163 exec_list_make_empty(&b->functions);
5164 b->entry_point_stage = stage;
5165 b->entry_point_name = entry_point_name;
5166 b->options = dup_options;
5167
5168 /*
5169 * Handle the SPIR-V header (first 5 dwords).
5170 * Can't use vtx_assert() as the setjmp(3) target isn't initialized yet.
5171 */
5172 if (word_count <= 5)
5173 goto fail;
5174
5175 if (words[0] != SpvMagicNumber) {
5176 vtn_err("words[0] was 0x%x, want 0x%x", words[0], SpvMagicNumber);
5177 goto fail;
5178 }
5179 if (words[1] < 0x10000) {
5180 vtn_err("words[1] was 0x%x, want >= 0x10000", words[1]);
5181 goto fail;
5182 }
5183
5184 uint16_t generator_id = words[2] >> 16;
5185 uint16_t generator_version = words[2];
5186
5187 /* The first GLSLang version bump actually 1.5 years after #179 was fixed
5188 * but this should at least let us shut the workaround off for modern
5189 * versions of GLSLang.
5190 */
5191 b->wa_glslang_179 = (generator_id == 8 && generator_version == 1);
5192
5193 /* In GLSLang commit 8297936dd6eb3, their handling of barrier() was fixed
5194 * to provide correct memory semantics on compute shader barrier()
5195 * commands. Prior to that, we need to fix them up ourselves. This
5196 * GLSLang fix caused them to bump to generator version 3.
5197 */
5198 b->wa_glslang_cs_barrier = (generator_id == 8 && generator_version < 3);
5199
5200 /* words[2] == generator magic */
5201 unsigned value_id_bound = words[3];
5202 if (words[4] != 0) {
5203 vtn_err("words[4] was %u, want 0", words[4]);
5204 goto fail;
5205 }
5206
5207 b->value_id_bound = value_id_bound;
5208 b->values = rzalloc_array(b, struct vtn_value, value_id_bound);
5209
5210 return b;
5211 fail:
5212 ralloc_free(b);
5213 return NULL;
5214 }
5215
5216 static nir_function *
5217 vtn_emit_kernel_entry_point_wrapper(struct vtn_builder *b,
5218 nir_function *entry_point)
5219 {
5220 vtn_assert(entry_point == b->entry_point->func->impl->function);
5221 vtn_fail_if(!entry_point->name, "entry points are required to have a name");
5222 const char *func_name =
5223 ralloc_asprintf(b->shader, "__wrapped_%s", entry_point->name);
5224
5225 /* we shouldn't have any inputs yet */
5226 vtn_assert(!entry_point->shader->num_inputs);
5227 vtn_assert(b->shader->info.stage == MESA_SHADER_KERNEL);
5228
5229 nir_function *main_entry_point = nir_function_create(b->shader, func_name);
5230 main_entry_point->impl = nir_function_impl_create(main_entry_point);
5231 nir_builder_init(&b->nb, main_entry_point->impl);
5232 b->nb.cursor = nir_after_cf_list(&main_entry_point->impl->body);
5233 b->func_param_idx = 0;
5234
5235 nir_call_instr *call = nir_call_instr_create(b->nb.shader, entry_point);
5236
5237 for (unsigned i = 0; i < entry_point->num_params; ++i) {
5238 struct vtn_type *param_type = b->entry_point->func->type->params[i];
5239
5240 /* consider all pointers to function memory to be parameters passed
5241 * by value
5242 */
5243 bool is_by_val = param_type->base_type == vtn_base_type_pointer &&
5244 param_type->storage_class == SpvStorageClassFunction;
5245
5246 /* input variable */
5247 nir_variable *in_var = rzalloc(b->nb.shader, nir_variable);
5248 in_var->data.mode = nir_var_shader_in;
5249 in_var->data.read_only = true;
5250 in_var->data.location = i;
5251
5252 if (is_by_val)
5253 in_var->type = param_type->deref->type;
5254 else
5255 in_var->type = param_type->type;
5256
5257 nir_shader_add_variable(b->nb.shader, in_var);
5258 b->nb.shader->num_inputs++;
5259
5260 /* we have to copy the entire variable into function memory */
5261 if (is_by_val) {
5262 nir_variable *copy_var =
5263 nir_local_variable_create(main_entry_point->impl, in_var->type,
5264 "copy_in");
5265 nir_copy_var(&b->nb, copy_var, in_var);
5266 call->params[i] =
5267 nir_src_for_ssa(&nir_build_deref_var(&b->nb, copy_var)->dest.ssa);
5268 } else {
5269 call->params[i] = nir_src_for_ssa(nir_load_var(&b->nb, in_var));
5270 }
5271 }
5272
5273 nir_builder_instr_insert(&b->nb, &call->instr);
5274
5275 return main_entry_point;
5276 }
5277
5278 nir_shader *
5279 spirv_to_nir(const uint32_t *words, size_t word_count,
5280 struct nir_spirv_specialization *spec, unsigned num_spec,
5281 gl_shader_stage stage, const char *entry_point_name,
5282 const struct spirv_to_nir_options *options,
5283 const nir_shader_compiler_options *nir_options)
5284
5285 {
5286 const uint32_t *word_end = words + word_count;
5287
5288 struct vtn_builder *b = vtn_create_builder(words, word_count,
5289 stage, entry_point_name,
5290 options);
5291
5292 if (b == NULL)
5293 return NULL;
5294
5295 /* See also _vtn_fail() */
5296 if (setjmp(b->fail_jump)) {
5297 ralloc_free(b);
5298 return NULL;
5299 }
5300
5301 /* Skip the SPIR-V header, handled at vtn_create_builder */
5302 words+= 5;
5303
5304 b->shader = nir_shader_create(b, stage, nir_options, NULL);
5305
5306 /* Handle all the preamble instructions */
5307 words = vtn_foreach_instruction(b, words, word_end,
5308 vtn_handle_preamble_instruction);
5309
5310 if (b->entry_point == NULL) {
5311 vtn_fail("Entry point not found");
5312 ralloc_free(b);
5313 return NULL;
5314 }
5315
5316 /* Set shader info defaults */
5317 if (stage == MESA_SHADER_GEOMETRY)
5318 b->shader->info.gs.invocations = 1;
5319
5320 /* Parse rounding mode execution modes. This has to happen earlier than
5321 * other changes in the execution modes since they can affect, for example,
5322 * the result of the floating point constants.
5323 */
5324 vtn_foreach_execution_mode(b, b->entry_point,
5325 vtn_handle_rounding_mode_in_execution_mode, NULL);
5326
5327 b->specializations = spec;
5328 b->num_specializations = num_spec;
5329
5330 /* Handle all variable, type, and constant instructions */
5331 words = vtn_foreach_instruction(b, words, word_end,
5332 vtn_handle_variable_or_type_instruction);
5333
5334 /* Parse execution modes */
5335 vtn_foreach_execution_mode(b, b->entry_point,
5336 vtn_handle_execution_mode, NULL);
5337
5338 if (b->workgroup_size_builtin) {
5339 vtn_assert(b->workgroup_size_builtin->type->type ==
5340 glsl_vector_type(GLSL_TYPE_UINT, 3));
5341
5342 nir_const_value *const_size =
5343 b->workgroup_size_builtin->constant->values;
5344
5345 b->shader->info.cs.local_size[0] = const_size[0].u32;
5346 b->shader->info.cs.local_size[1] = const_size[1].u32;
5347 b->shader->info.cs.local_size[2] = const_size[2].u32;
5348 }
5349
5350 /* Set types on all vtn_values */
5351 vtn_foreach_instruction(b, words, word_end, vtn_set_instruction_result_type);
5352
5353 vtn_build_cfg(b, words, word_end);
5354
5355 assert(b->entry_point->value_type == vtn_value_type_function);
5356 b->entry_point->func->referenced = true;
5357
5358 bool progress;
5359 do {
5360 progress = false;
5361 foreach_list_typed(struct vtn_function, func, node, &b->functions) {
5362 if (func->referenced && !func->emitted) {
5363 b->const_table = _mesa_pointer_hash_table_create(b);
5364
5365 vtn_function_emit(b, func, vtn_handle_body_instruction);
5366 progress = true;
5367 }
5368 }
5369 } while (progress);
5370
5371 vtn_assert(b->entry_point->value_type == vtn_value_type_function);
5372 nir_function *entry_point = b->entry_point->func->impl->function;
5373 vtn_assert(entry_point);
5374
5375 /* post process entry_points with input params */
5376 if (entry_point->num_params && b->shader->info.stage == MESA_SHADER_KERNEL)
5377 entry_point = vtn_emit_kernel_entry_point_wrapper(b, entry_point);
5378
5379 entry_point->is_entrypoint = true;
5380
5381 /* When multiple shader stages exist in the same SPIR-V module, we
5382 * generate input and output variables for every stage, in the same
5383 * NIR program. These dead variables can be invalid NIR. For example,
5384 * TCS outputs must be per-vertex arrays (or decorated 'patch'), while
5385 * VS output variables wouldn't be.
5386 *
5387 * To ensure we have valid NIR, we eliminate any dead inputs and outputs
5388 * right away. In order to do so, we must lower any constant initializers
5389 * on outputs so nir_remove_dead_variables sees that they're written to.
5390 */
5391 nir_lower_constant_initializers(b->shader, nir_var_shader_out);
5392 nir_remove_dead_variables(b->shader,
5393 nir_var_shader_in | nir_var_shader_out);
5394
5395 /* We sometimes generate bogus derefs that, while never used, give the
5396 * validator a bit of heartburn. Run dead code to get rid of them.
5397 */
5398 nir_opt_dce(b->shader);
5399
5400 /* Unparent the shader from the vtn_builder before we delete the builder */
5401 ralloc_steal(NULL, b->shader);
5402
5403 nir_shader *shader = b->shader;
5404 ralloc_free(b);
5405
5406 return shader;
5407 }