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