projects / mesa.git / blob
? search:
summary | shortlog | log | commit | commitdiff | tree
history | raw | HEAD
spirv: Make push constants an offset-based pointer
[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 <stdio.h>
36
37 void
38 vtn_log(struct vtn_builder *b, enum nir_spirv_debug_level level,
39 size_t spirv_offset, const char *message)
40 {
41 if (b->options->debug.func) {
42 b->options->debug.func(b->options->debug.private_data,
43 level, spirv_offset, message);
44 }
45
46 #ifndef NDEBUG
47 if (level >= NIR_SPIRV_DEBUG_LEVEL_WARNING)
48 fprintf(stderr, "%s\n", message);
49 #endif
50 }
51
52 void
53 vtn_logf(struct vtn_builder *b, enum nir_spirv_debug_level level,
54 size_t spirv_offset, const char *fmt, ...)
55 {
56 va_list args;
57 char *msg;
58
59 va_start(args, fmt);
60 msg = ralloc_vasprintf(NULL, fmt, args);
61 va_end(args);
62
63 vtn_log(b, level, spirv_offset, msg);
64
65 ralloc_free(msg);
66 }
67
68 static void
69 vtn_log_err(struct vtn_builder *b,
70 enum nir_spirv_debug_level level, const char *prefix,
71 const char *file, unsigned line,
72 const char *fmt, va_list args)
73 {
74 char *msg;
75
76 msg = ralloc_strdup(NULL, prefix);
77
78 #ifndef NDEBUG
79 ralloc_asprintf_append(&msg, " In file %s:%u\n", file, line);
80 #endif
81
82 ralloc_asprintf_append(&msg, " ");
83
84 ralloc_vasprintf_append(&msg, fmt, args);
85
86 ralloc_asprintf_append(&msg, "\n %zu bytes into the SPIR-V binary",
87 b->spirv_offset);
88
89 if (b->file) {
90 ralloc_asprintf_append(&msg,
91 "\n in SPIR-V source file %s, line %d, col %d",
92 b->file, b->line, b->col);
93 }
94
95 vtn_log(b, level, b->spirv_offset, msg);
96
97 ralloc_free(msg);
98 }
99
100 static void
101 vtn_dump_shader(struct vtn_builder *b, const char *path, const char *prefix)
102 {
103 static int idx = 0;
104
105 char filename[1024];
106 int len = snprintf(filename, sizeof(filename), "%s/%s-%d.spirv",
107 path, prefix, idx++);
108 if (len < 0 || len >= sizeof(filename))
109 return;
110
111 FILE *f = fopen(filename, "w");
112 if (f == NULL)
113 return;
114
115 fwrite(b->spirv, sizeof(*b->spirv), b->spirv_word_count, f);
116 fclose(f);
117
118 vtn_info("SPIR-V shader dumped to %s", filename);
119 }
120
121 void
122 _vtn_warn(struct vtn_builder *b, const char *file, unsigned line,
123 const char *fmt, ...)
124 {
125 va_list args;
126
127 va_start(args, fmt);
128 vtn_log_err(b, NIR_SPIRV_DEBUG_LEVEL_WARNING, "SPIR-V WARNING:\n",
129 file, line, fmt, args);
130 va_end(args);
131 }
132
133 void
134 _vtn_err(struct vtn_builder *b, const char *file, unsigned line,
135 const char *fmt, ...)
136 {
137 va_list args;
138
139 va_start(args, fmt);
140 vtn_log_err(b, NIR_SPIRV_DEBUG_LEVEL_ERROR, "SPIR-V ERROR:\n",
141 file, line, fmt, args);
142 va_end(args);
143 }
144
145 void
146 _vtn_fail(struct vtn_builder *b, const char *file, unsigned line,
147 const char *fmt, ...)
148 {
149 va_list args;
150
151 va_start(args, fmt);
152 vtn_log_err(b, NIR_SPIRV_DEBUG_LEVEL_ERROR, "SPIR-V parsing FAILED:\n",
153 file, line, fmt, args);
154 va_end(args);
155
156 const char *dump_path = getenv("MESA_SPIRV_FAIL_DUMP_PATH");
157 if (dump_path)
158 vtn_dump_shader(b, dump_path, "fail");
159
160 longjmp(b->fail_jump, 1);
161 }
162
163 struct spec_constant_value {
164 bool is_double;
165 union {
166 uint32_t data32;
167 uint64_t data64;
168 };
169 };
170
171 static struct vtn_ssa_value *
172 vtn_undef_ssa_value(struct vtn_builder *b, const struct glsl_type *type)
173 {
174 struct vtn_ssa_value *val = rzalloc(b, struct vtn_ssa_value);
175 val->type = type;
176
177 if (glsl_type_is_vector_or_scalar(type)) {
178 unsigned num_components = glsl_get_vector_elements(val->type);
179 unsigned bit_size = glsl_get_bit_size(val->type);
180 val->def = nir_ssa_undef(&b->nb, num_components, bit_size);
181 } else {
182 unsigned elems = glsl_get_length(val->type);
183 val->elems = ralloc_array(b, struct vtn_ssa_value *, elems);
184 if (glsl_type_is_matrix(type)) {
185 const struct glsl_type *elem_type =
186 glsl_vector_type(glsl_get_base_type(type),
187 glsl_get_vector_elements(type));
188
189 for (unsigned i = 0; i < elems; i++)
190 val->elems[i] = vtn_undef_ssa_value(b, elem_type);
191 } else if (glsl_type_is_array(type)) {
192 const struct glsl_type *elem_type = glsl_get_array_element(type);
193 for (unsigned i = 0; i < elems; i++)
194 val->elems[i] = vtn_undef_ssa_value(b, elem_type);
195 } else {
196 for (unsigned i = 0; i < elems; i++) {
197 const struct glsl_type *elem_type = glsl_get_struct_field(type, i);
198 val->elems[i] = vtn_undef_ssa_value(b, elem_type);
199 }
200 }
201 }
202
203 return val;
204 }
205
206 static struct vtn_ssa_value *
207 vtn_const_ssa_value(struct vtn_builder *b, nir_constant *constant,
208 const struct glsl_type *type)
209 {
210 struct hash_entry *entry = _mesa_hash_table_search(b->const_table, constant);
211
212 if (entry)
213 return entry->data;
214
215 struct vtn_ssa_value *val = rzalloc(b, struct vtn_ssa_value);
216 val->type = type;
217
218 switch (glsl_get_base_type(type)) {
219 case GLSL_TYPE_INT:
220 case GLSL_TYPE_UINT:
221 case GLSL_TYPE_INT16:
222 case GLSL_TYPE_UINT16:
223 case GLSL_TYPE_UINT8:
224 case GLSL_TYPE_INT8:
225 case GLSL_TYPE_INT64:
226 case GLSL_TYPE_UINT64:
227 case GLSL_TYPE_BOOL:
228 case GLSL_TYPE_FLOAT:
229 case GLSL_TYPE_FLOAT16:
230 case GLSL_TYPE_DOUBLE: {
231 int bit_size = glsl_get_bit_size(type);
232 if (glsl_type_is_vector_or_scalar(type)) {
233 unsigned num_components = glsl_get_vector_elements(val->type);
234 nir_load_const_instr *load =
235 nir_load_const_instr_create(b->shader, num_components, bit_size);
236
237 load->value = constant->values[0];
238
239 nir_instr_insert_before_cf_list(&b->nb.impl->body, &load->instr);
240 val->def = &load->def;
241 } else {
242 assert(glsl_type_is_matrix(type));
243 unsigned rows = glsl_get_vector_elements(val->type);
244 unsigned columns = glsl_get_matrix_columns(val->type);
245 val->elems = ralloc_array(b, struct vtn_ssa_value *, columns);
246
247 for (unsigned i = 0; i < columns; i++) {
248 struct vtn_ssa_value *col_val = rzalloc(b, struct vtn_ssa_value);
249 col_val->type = glsl_get_column_type(val->type);
250 nir_load_const_instr *load =
251 nir_load_const_instr_create(b->shader, rows, bit_size);
252
253 load->value = constant->values[i];
254
255 nir_instr_insert_before_cf_list(&b->nb.impl->body, &load->instr);
256 col_val->def = &load->def;
257
258 val->elems[i] = col_val;
259 }
260 }
261 break;
262 }
263
264 case GLSL_TYPE_ARRAY: {
265 unsigned elems = glsl_get_length(val->type);
266 val->elems = ralloc_array(b, struct vtn_ssa_value *, elems);
267 const struct glsl_type *elem_type = glsl_get_array_element(val->type);
268 for (unsigned i = 0; i < elems; i++)
269 val->elems[i] = vtn_const_ssa_value(b, constant->elements[i],
270 elem_type);
271 break;
272 }
273
274 case GLSL_TYPE_STRUCT: {
275 unsigned elems = glsl_get_length(val->type);
276 val->elems = ralloc_array(b, struct vtn_ssa_value *, elems);
277 for (unsigned i = 0; i < elems; i++) {
278 const struct glsl_type *elem_type =
279 glsl_get_struct_field(val->type, i);
280 val->elems[i] = vtn_const_ssa_value(b, constant->elements[i],
281 elem_type);
282 }
283 break;
284 }
285
286 default:
287 vtn_fail("bad constant type");
288 }
289
290 return val;
291 }
292
293 struct vtn_ssa_value *
294 vtn_ssa_value(struct vtn_builder *b, uint32_t value_id)
295 {
296 struct vtn_value *val = vtn_untyped_value(b, value_id);
297 switch (val->value_type) {
298 case vtn_value_type_undef:
299 return vtn_undef_ssa_value(b, val->type->type);
300
301 case vtn_value_type_constant:
302 return vtn_const_ssa_value(b, val->constant, val->type->type);
303
304 case vtn_value_type_ssa:
305 return val->ssa;
306
307 case vtn_value_type_pointer:
308 vtn_assert(val->pointer->ptr_type && val->pointer->ptr_type->type);
309 struct vtn_ssa_value *ssa =
310 vtn_create_ssa_value(b, val->pointer->ptr_type->type);
311 ssa->def = vtn_pointer_to_ssa(b, val->pointer);
312 return ssa;
313
314 default:
315 vtn_fail("Invalid type for an SSA value");
316 }
317 }
318
319 static char *
320 vtn_string_literal(struct vtn_builder *b, const uint32_t *words,
321 unsigned word_count, unsigned *words_used)
322 {
323 char *dup = ralloc_strndup(b, (char *)words, word_count * sizeof(*words));
324 if (words_used) {
325 /* Ammount of space taken by the string (including the null) */
326 unsigned len = strlen(dup) + 1;
327 *words_used = DIV_ROUND_UP(len, sizeof(*words));
328 }
329 return dup;
330 }
331
332 const uint32_t *
333 vtn_foreach_instruction(struct vtn_builder *b, const uint32_t *start,
334 const uint32_t *end, vtn_instruction_handler handler)
335 {
336 b->file = NULL;
337 b->line = -1;
338 b->col = -1;
339
340 const uint32_t *w = start;
341 while (w < end) {
342 SpvOp opcode = w[0] & SpvOpCodeMask;
343 unsigned count = w[0] >> SpvWordCountShift;
344 vtn_assert(count >= 1 && w + count <= end);
345
346 b->spirv_offset = (uint8_t *)w - (uint8_t *)b->spirv;
347
348 switch (opcode) {
349 case SpvOpNop:
350 break; /* Do nothing */
351
352 case SpvOpLine:
353 b->file = vtn_value(b, w[1], vtn_value_type_string)->str;
354 b->line = w[2];
355 b->col = w[3];
356 break;
357
358 case SpvOpNoLine:
359 b->file = NULL;
360 b->line = -1;
361 b->col = -1;
362 break;
363
364 default:
365 if (!handler(b, opcode, w, count))
366 return w;
367 break;
368 }
369
370 w += count;
371 }
372
373 b->spirv_offset = 0;
374 b->file = NULL;
375 b->line = -1;
376 b->col = -1;
377
378 assert(w == end);
379 return w;
380 }
381
382 static void
383 vtn_handle_extension(struct vtn_builder *b, SpvOp opcode,
384 const uint32_t *w, unsigned count)
385 {
386 switch (opcode) {
387 case SpvOpExtInstImport: {
388 struct vtn_value *val = vtn_push_value(b, w[1], vtn_value_type_extension);
389 if (strcmp((const char *)&w[2], "GLSL.std.450") == 0) {
390 val->ext_handler = vtn_handle_glsl450_instruction;
391 } else if ((strcmp((const char *)&w[2], "SPV_AMD_gcn_shader") == 0)
392 && (b->options && b->options->caps.gcn_shader)) {
393 val->ext_handler = vtn_handle_amd_gcn_shader_instruction;
394 } else if ((strcmp((const char *)&w[2], "SPV_AMD_shader_trinary_minmax") == 0)
395 && (b->options && b->options->caps.trinary_minmax)) {
396 val->ext_handler = vtn_handle_amd_shader_trinary_minmax_instruction;
397 } else {
398 vtn_fail("Unsupported extension");
399 }
400 break;
401 }
402
403 case SpvOpExtInst: {
404 struct vtn_value *val = vtn_value(b, w[3], vtn_value_type_extension);
405 bool handled = val->ext_handler(b, w[4], w, count);
406 vtn_assert(handled);
407 break;
408 }
409
410 default:
411 vtn_fail("Unhandled opcode");
412 }
413 }
414
415 static void
416 _foreach_decoration_helper(struct vtn_builder *b,
417 struct vtn_value *base_value,
418 int parent_member,
419 struct vtn_value *value,
420 vtn_decoration_foreach_cb cb, void *data)
421 {
422 for (struct vtn_decoration *dec = value->decoration; dec; dec = dec->next) {
423 int member;
424 if (dec->scope == VTN_DEC_DECORATION) {
425 member = parent_member;
426 } else if (dec->scope >= VTN_DEC_STRUCT_MEMBER0) {
427 vtn_fail_if(value->value_type != vtn_value_type_type ||
428 value->type->base_type != vtn_base_type_struct,
429 "OpMemberDecorate and OpGroupMemberDecorate are only "
430 "allowed on OpTypeStruct");
431 /* This means we haven't recursed yet */
432 assert(value == base_value);
433
434 member = dec->scope - VTN_DEC_STRUCT_MEMBER0;
435
436 vtn_fail_if(member >= base_value->type->length,
437 "OpMemberDecorate specifies member %d but the "
438 "OpTypeStruct has only %u members",
439 member, base_value->type->length);
440 } else {
441 /* Not a decoration */
442 assert(dec->scope == VTN_DEC_EXECUTION_MODE);
443 continue;
444 }
445
446 if (dec->group) {
447 assert(dec->group->value_type == vtn_value_type_decoration_group);
448 _foreach_decoration_helper(b, base_value, member, dec->group,
449 cb, data);
450 } else {
451 cb(b, base_value, member, dec, data);
452 }
453 }
454 }
455
456 /** Iterates (recursively if needed) over all of the decorations on a value
457 *
458 * This function iterates over all of the decorations applied to a given
459 * value. If it encounters a decoration group, it recurses into the group
460 * and iterates over all of those decorations as well.
461 */
462 void
463 vtn_foreach_decoration(struct vtn_builder *b, struct vtn_value *value,
464 vtn_decoration_foreach_cb cb, void *data)
465 {
466 _foreach_decoration_helper(b, value, -1, value, cb, data);
467 }
468
469 void
470 vtn_foreach_execution_mode(struct vtn_builder *b, struct vtn_value *value,
471 vtn_execution_mode_foreach_cb cb, void *data)
472 {
473 for (struct vtn_decoration *dec = value->decoration; dec; dec = dec->next) {
474 if (dec->scope != VTN_DEC_EXECUTION_MODE)
475 continue;
476
477 assert(dec->group == NULL);
478 cb(b, value, dec, data);
479 }
480 }
481
482 void
483 vtn_handle_decoration(struct vtn_builder *b, SpvOp opcode,
484 const uint32_t *w, unsigned count)
485 {
486 const uint32_t *w_end = w + count;
487 const uint32_t target = w[1];
488 w += 2;
489
490 switch (opcode) {
491 case SpvOpDecorationGroup:
492 vtn_push_value(b, target, vtn_value_type_decoration_group);
493 break;
494
495 case SpvOpDecorate:
496 case SpvOpMemberDecorate:
497 case SpvOpExecutionMode: {
498 struct vtn_value *val = vtn_untyped_value(b, target);
499
500 struct vtn_decoration *dec = rzalloc(b, struct vtn_decoration);
501 switch (opcode) {
502 case SpvOpDecorate:
503 dec->scope = VTN_DEC_DECORATION;
504 break;
505 case SpvOpMemberDecorate:
506 dec->scope = VTN_DEC_STRUCT_MEMBER0 + *(w++);
507 vtn_fail_if(dec->scope < VTN_DEC_STRUCT_MEMBER0, /* overflow */
508 "Member argument of OpMemberDecorate too large");
509 break;
510 case SpvOpExecutionMode:
511 dec->scope = VTN_DEC_EXECUTION_MODE;
512 break;
513 default:
514 unreachable("Invalid decoration opcode");
515 }
516 dec->decoration = *(w++);
517 dec->literals = w;
518
519 /* Link into the list */
520 dec->next = val->decoration;
521 val->decoration = dec;
522 break;
523 }
524
525 case SpvOpGroupMemberDecorate:
526 case SpvOpGroupDecorate: {
527 struct vtn_value *group =
528 vtn_value(b, target, vtn_value_type_decoration_group);
529
530 for (; w < w_end; w++) {
531 struct vtn_value *val = vtn_untyped_value(b, *w);
532 struct vtn_decoration *dec = rzalloc(b, struct vtn_decoration);
533
534 dec->group = group;
535 if (opcode == SpvOpGroupDecorate) {
536 dec->scope = VTN_DEC_DECORATION;
537 } else {
538 dec->scope = VTN_DEC_STRUCT_MEMBER0 + *(++w);
539 vtn_fail_if(dec->scope < 0, /* Check for overflow */
540 "Member argument of OpGroupMemberDecorate too large");
541 }
542
543 /* Link into the list */
544 dec->next = val->decoration;
545 val->decoration = dec;
546 }
547 break;
548 }
549
550 default:
551 unreachable("Unhandled opcode");
552 }
553 }
554
555 struct member_decoration_ctx {
556 unsigned num_fields;
557 struct glsl_struct_field *fields;
558 struct vtn_type *type;
559 };
560
561 /** Returns true if two types are "compatible", i.e. you can do an OpLoad,
562 * OpStore, or OpCopyMemory between them without breaking anything.
563 * Technically, the SPIR-V rules require the exact same type ID but this lets
564 * us internally be a bit looser.
565 */
566 bool
567 vtn_types_compatible(struct vtn_builder *b,
568 struct vtn_type *t1, struct vtn_type *t2)
569 {
570 if (t1->id == t2->id)
571 return true;
572
573 if (t1->base_type != t2->base_type)
574 return false;
575
576 switch (t1->base_type) {
577 case vtn_base_type_void:
578 case vtn_base_type_scalar:
579 case vtn_base_type_vector:
580 case vtn_base_type_matrix:
581 case vtn_base_type_image:
582 case vtn_base_type_sampler:
583 case vtn_base_type_sampled_image:
584 return t1->type == t2->type;
585
586 case vtn_base_type_array:
587 return t1->length == t2->length &&
588 vtn_types_compatible(b, t1->array_element, t2->array_element);
589
590 case vtn_base_type_pointer:
591 return vtn_types_compatible(b, t1->deref, t2->deref);
592
593 case vtn_base_type_struct:
594 if (t1->length != t2->length)
595 return false;
596
597 for (unsigned i = 0; i < t1->length; i++) {
598 if (!vtn_types_compatible(b, t1->members[i], t2->members[i]))
599 return false;
600 }
601 return true;
602
603 case vtn_base_type_function:
604 /* This case shouldn't get hit since you can't copy around function
605 * types. Just require them to be identical.
606 */
607 return false;
608 }
609
610 vtn_fail("Invalid base type");
611 }
612
613 /* does a shallow copy of a vtn_type */
614
615 static struct vtn_type *
616 vtn_type_copy(struct vtn_builder *b, struct vtn_type *src)
617 {
618 struct vtn_type *dest = ralloc(b, struct vtn_type);
619 *dest = *src;
620
621 switch (src->base_type) {
622 case vtn_base_type_void:
623 case vtn_base_type_scalar:
624 case vtn_base_type_vector:
625 case vtn_base_type_matrix:
626 case vtn_base_type_array:
627 case vtn_base_type_pointer:
628 case vtn_base_type_image:
629 case vtn_base_type_sampler:
630 case vtn_base_type_sampled_image:
631 /* Nothing more to do */
632 break;
633
634 case vtn_base_type_struct:
635 dest->members = ralloc_array(b, struct vtn_type *, src->length);
636 memcpy(dest->members, src->members,
637 src->length * sizeof(src->members[0]));
638
639 dest->offsets = ralloc_array(b, unsigned, src->length);
640 memcpy(dest->offsets, src->offsets,
641 src->length * sizeof(src->offsets[0]));
642 break;
643
644 case vtn_base_type_function:
645 dest->params = ralloc_array(b, struct vtn_type *, src->length);
646 memcpy(dest->params, src->params, src->length * sizeof(src->params[0]));
647 break;
648 }
649
650 return dest;
651 }
652
653 static struct vtn_type *
654 mutable_matrix_member(struct vtn_builder *b, struct vtn_type *type, int member)
655 {
656 type->members[member] = vtn_type_copy(b, type->members[member]);
657 type = type->members[member];
658
659 /* We may have an array of matrices.... Oh, joy! */
660 while (glsl_type_is_array(type->type)) {
661 type->array_element = vtn_type_copy(b, type->array_element);
662 type = type->array_element;
663 }
664
665 vtn_assert(glsl_type_is_matrix(type->type));
666
667 return type;
668 }
669
670 static void
671 struct_member_decoration_cb(struct vtn_builder *b,
672 struct vtn_value *val, int member,
673 const struct vtn_decoration *dec, void *void_ctx)
674 {
675 struct member_decoration_ctx *ctx = void_ctx;
676
677 if (member < 0)
678 return;
679
680 assert(member < ctx->num_fields);
681
682 switch (dec->decoration) {
683 case SpvDecorationNonWritable:
684 case SpvDecorationNonReadable:
685 case SpvDecorationRelaxedPrecision:
686 case SpvDecorationVolatile:
687 case SpvDecorationCoherent:
688 case SpvDecorationUniform:
689 break; /* FIXME: Do nothing with this for now. */
690 case SpvDecorationNoPerspective:
691 ctx->fields[member].interpolation = INTERP_MODE_NOPERSPECTIVE;
692 break;
693 case SpvDecorationFlat:
694 ctx->fields[member].interpolation = INTERP_MODE_FLAT;
695 break;
696 case SpvDecorationCentroid:
697 ctx->fields[member].centroid = true;
698 break;
699 case SpvDecorationSample:
700 ctx->fields[member].sample = true;
701 break;
702 case SpvDecorationStream:
703 /* Vulkan only allows one GS stream */
704 vtn_assert(dec->literals[0] == 0);
705 break;
706 case SpvDecorationLocation:
707 ctx->fields[member].location = dec->literals[0];
708 break;
709 case SpvDecorationComponent:
710 break; /* FIXME: What should we do with these? */
711 case SpvDecorationBuiltIn:
712 ctx->type->members[member] = vtn_type_copy(b, ctx->type->members[member]);
713 ctx->type->members[member]->is_builtin = true;
714 ctx->type->members[member]->builtin = dec->literals[0];
715 ctx->type->builtin_block = true;
716 break;
717 case SpvDecorationOffset:
718 ctx->type->offsets[member] = dec->literals[0];
719 break;
720 case SpvDecorationMatrixStride:
721 /* Handled as a second pass */
722 break;
723 case SpvDecorationColMajor:
724 break; /* Nothing to do here. Column-major is the default. */
725 case SpvDecorationRowMajor:
726 mutable_matrix_member(b, ctx->type, member)->row_major = true;
727 break;
728
729 case SpvDecorationPatch:
730 break;
731
732 case SpvDecorationSpecId:
733 case SpvDecorationBlock:
734 case SpvDecorationBufferBlock:
735 case SpvDecorationArrayStride:
736 case SpvDecorationGLSLShared:
737 case SpvDecorationGLSLPacked:
738 case SpvDecorationInvariant:
739 case SpvDecorationRestrict:
740 case SpvDecorationAliased:
741 case SpvDecorationConstant:
742 case SpvDecorationIndex:
743 case SpvDecorationBinding:
744 case SpvDecorationDescriptorSet:
745 case SpvDecorationLinkageAttributes:
746 case SpvDecorationNoContraction:
747 case SpvDecorationInputAttachmentIndex:
748 vtn_warn("Decoration not allowed on struct members: %s",
749 spirv_decoration_to_string(dec->decoration));
750 break;
751
752 case SpvDecorationXfbBuffer:
753 case SpvDecorationXfbStride:
754 vtn_warn("Vulkan does not have transform feedback");
755 break;
756
757 case SpvDecorationCPacked:
758 case SpvDecorationSaturatedConversion:
759 case SpvDecorationFuncParamAttr:
760 case SpvDecorationFPRoundingMode:
761 case SpvDecorationFPFastMathMode:
762 case SpvDecorationAlignment:
763 vtn_warn("Decoration only allowed for CL-style kernels: %s",
764 spirv_decoration_to_string(dec->decoration));
765 break;
766
767 default:
768 vtn_fail("Unhandled decoration");
769 }
770 }
771
772 /* Matrix strides are handled as a separate pass because we need to know
773 * whether the matrix is row-major or not first.
774 */
775 static void
776 struct_member_matrix_stride_cb(struct vtn_builder *b,
777 struct vtn_value *val, int member,
778 const struct vtn_decoration *dec,
779 void *void_ctx)
780 {
781 if (dec->decoration != SpvDecorationMatrixStride)
782 return;
783
784 vtn_fail_if(member < 0,
785 "The MatrixStride decoration is only allowed on members "
786 "of OpTypeStruct");
787
788 struct member_decoration_ctx *ctx = void_ctx;
789
790 struct vtn_type *mat_type = mutable_matrix_member(b, ctx->type, member);
791 if (mat_type->row_major) {
792 mat_type->array_element = vtn_type_copy(b, mat_type->array_element);
793 mat_type->stride = mat_type->array_element->stride;
794 mat_type->array_element->stride = dec->literals[0];
795 } else {
796 vtn_assert(mat_type->array_element->stride > 0);
797 mat_type->stride = dec->literals[0];
798 }
799 }
800
801 static void
802 type_decoration_cb(struct vtn_builder *b,
803 struct vtn_value *val, int member,
804 const struct vtn_decoration *dec, void *ctx)
805 {
806 struct vtn_type *type = val->type;
807
808 if (member != -1) {
809 /* This should have been handled by OpTypeStruct */
810 assert(val->type->base_type == vtn_base_type_struct);
811 assert(member >= 0 && member < val->type->length);
812 return;
813 }
814
815 switch (dec->decoration) {
816 case SpvDecorationArrayStride:
817 vtn_assert(type->base_type == vtn_base_type_matrix ||
818 type->base_type == vtn_base_type_array ||
819 type->base_type == vtn_base_type_pointer);
820 type->stride = dec->literals[0];
821 break;
822 case SpvDecorationBlock:
823 vtn_assert(type->base_type == vtn_base_type_struct);
824 type->block = true;
825 break;
826 case SpvDecorationBufferBlock:
827 vtn_assert(type->base_type == vtn_base_type_struct);
828 type->buffer_block = true;
829 break;
830 case SpvDecorationGLSLShared:
831 case SpvDecorationGLSLPacked:
832 /* Ignore these, since we get explicit offsets anyways */
833 break;
834
835 case SpvDecorationRowMajor:
836 case SpvDecorationColMajor:
837 case SpvDecorationMatrixStride:
838 case SpvDecorationBuiltIn:
839 case SpvDecorationNoPerspective:
840 case SpvDecorationFlat:
841 case SpvDecorationPatch:
842 case SpvDecorationCentroid:
843 case SpvDecorationSample:
844 case SpvDecorationVolatile:
845 case SpvDecorationCoherent:
846 case SpvDecorationNonWritable:
847 case SpvDecorationNonReadable:
848 case SpvDecorationUniform:
849 case SpvDecorationStream:
850 case SpvDecorationLocation:
851 case SpvDecorationComponent:
852 case SpvDecorationOffset:
853 case SpvDecorationXfbBuffer:
854 case SpvDecorationXfbStride:
855 vtn_warn("Decoration only allowed for struct members: %s",
856 spirv_decoration_to_string(dec->decoration));
857 break;
858
859 case SpvDecorationRelaxedPrecision:
860 case SpvDecorationSpecId:
861 case SpvDecorationInvariant:
862 case SpvDecorationRestrict:
863 case SpvDecorationAliased:
864 case SpvDecorationConstant:
865 case SpvDecorationIndex:
866 case SpvDecorationBinding:
867 case SpvDecorationDescriptorSet:
868 case SpvDecorationLinkageAttributes:
869 case SpvDecorationNoContraction:
870 case SpvDecorationInputAttachmentIndex:
871 vtn_warn("Decoration not allowed on types: %s",
872 spirv_decoration_to_string(dec->decoration));
873 break;
874
875 case SpvDecorationCPacked:
876 case SpvDecorationSaturatedConversion:
877 case SpvDecorationFuncParamAttr:
878 case SpvDecorationFPRoundingMode:
879 case SpvDecorationFPFastMathMode:
880 case SpvDecorationAlignment:
881 vtn_warn("Decoration only allowed for CL-style kernels: %s",
882 spirv_decoration_to_string(dec->decoration));
883 break;
884
885 default:
886 vtn_fail("Unhandled decoration");
887 }
888 }
889
890 static unsigned
891 translate_image_format(struct vtn_builder *b, SpvImageFormat format)
892 {
893 switch (format) {
894 case SpvImageFormatUnknown: return 0; /* GL_NONE */
895 case SpvImageFormatRgba32f: return 0x8814; /* GL_RGBA32F */
896 case SpvImageFormatRgba16f: return 0x881A; /* GL_RGBA16F */
897 case SpvImageFormatR32f: return 0x822E; /* GL_R32F */
898 case SpvImageFormatRgba8: return 0x8058; /* GL_RGBA8 */
899 case SpvImageFormatRgba8Snorm: return 0x8F97; /* GL_RGBA8_SNORM */
900 case SpvImageFormatRg32f: return 0x8230; /* GL_RG32F */
901 case SpvImageFormatRg16f: return 0x822F; /* GL_RG16F */
902 case SpvImageFormatR11fG11fB10f: return 0x8C3A; /* GL_R11F_G11F_B10F */
903 case SpvImageFormatR16f: return 0x822D; /* GL_R16F */
904 case SpvImageFormatRgba16: return 0x805B; /* GL_RGBA16 */
905 case SpvImageFormatRgb10A2: return 0x8059; /* GL_RGB10_A2 */
906 case SpvImageFormatRg16: return 0x822C; /* GL_RG16 */
907 case SpvImageFormatRg8: return 0x822B; /* GL_RG8 */
908 case SpvImageFormatR16: return 0x822A; /* GL_R16 */
909 case SpvImageFormatR8: return 0x8229; /* GL_R8 */
910 case SpvImageFormatRgba16Snorm: return 0x8F9B; /* GL_RGBA16_SNORM */
911 case SpvImageFormatRg16Snorm: return 0x8F99; /* GL_RG16_SNORM */
912 case SpvImageFormatRg8Snorm: return 0x8F95; /* GL_RG8_SNORM */
913 case SpvImageFormatR16Snorm: return 0x8F98; /* GL_R16_SNORM */
914 case SpvImageFormatR8Snorm: return 0x8F94; /* GL_R8_SNORM */
915 case SpvImageFormatRgba32i: return 0x8D82; /* GL_RGBA32I */
916 case SpvImageFormatRgba16i: return 0x8D88; /* GL_RGBA16I */
917 case SpvImageFormatRgba8i: return 0x8D8E; /* GL_RGBA8I */
918 case SpvImageFormatR32i: return 0x8235; /* GL_R32I */
919 case SpvImageFormatRg32i: return 0x823B; /* GL_RG32I */
920 case SpvImageFormatRg16i: return 0x8239; /* GL_RG16I */
921 case SpvImageFormatRg8i: return 0x8237; /* GL_RG8I */
922 case SpvImageFormatR16i: return 0x8233; /* GL_R16I */
923 case SpvImageFormatR8i: return 0x8231; /* GL_R8I */
924 case SpvImageFormatRgba32ui: return 0x8D70; /* GL_RGBA32UI */
925 case SpvImageFormatRgba16ui: return 0x8D76; /* GL_RGBA16UI */
926 case SpvImageFormatRgba8ui: return 0x8D7C; /* GL_RGBA8UI */
927 case SpvImageFormatR32ui: return 0x8236; /* GL_R32UI */
928 case SpvImageFormatRgb10a2ui: return 0x906F; /* GL_RGB10_A2UI */
929 case SpvImageFormatRg32ui: return 0x823C; /* GL_RG32UI */
930 case SpvImageFormatRg16ui: return 0x823A; /* GL_RG16UI */
931 case SpvImageFormatRg8ui: return 0x8238; /* GL_RG8UI */
932 case SpvImageFormatR16ui: return 0x8234; /* GL_R16UI */
933 case SpvImageFormatR8ui: return 0x8232; /* GL_R8UI */
934 default:
935 vtn_fail("Invalid image format");
936 }
937 }
938
939 static struct vtn_type *
940 vtn_type_layout_std430(struct vtn_builder *b, struct vtn_type *type,
941 uint32_t *size_out, uint32_t *align_out)
942 {
943 switch (type->base_type) {
944 case vtn_base_type_scalar: {
945 uint32_t comp_size = glsl_get_bit_size(type->type) / 8;
946 *size_out = comp_size;
947 *align_out = comp_size;
948 return type;
949 }
950
951 case vtn_base_type_vector: {
952 uint32_t comp_size = glsl_get_bit_size(type->type) / 8;
953 unsigned align_comps = type->length == 3 ? 4 : type->length;
954 *size_out = comp_size * type->length,
955 *align_out = comp_size * align_comps;
956 return type;
957 }
958
959 case vtn_base_type_matrix:
960 case vtn_base_type_array: {
961 /* We're going to add an array stride */
962 type = vtn_type_copy(b, type);
963 uint32_t elem_size, elem_align;
964 type->array_element = vtn_type_layout_std430(b, type->array_element,
965 &elem_size, &elem_align);
966 type->stride = vtn_align_u32(elem_size, elem_align);
967 *size_out = type->stride * type->length;
968 *align_out = elem_align;
969 return type;
970 }
971
972 case vtn_base_type_struct: {
973 /* We're going to add member offsets */
974 type = vtn_type_copy(b, type);
975 uint32_t offset = 0;
976 uint32_t align = 0;
977 for (unsigned i = 0; i < type->length; i++) {
978 uint32_t mem_size, mem_align;
979 type->members[i] = vtn_type_layout_std430(b, type->members[i],
980 &mem_size, &mem_align);
981 offset = vtn_align_u32(offset, mem_align);
982 type->offsets[i] = offset;
983 offset += mem_size;
984 align = MAX2(align, mem_align);
985 }
986 *size_out = offset;
987 *align_out = align;
988 return type;
989 }
990
991 default:
992 unreachable("Invalid SPIR-V type for std430");
993 }
994 }
995
996 static void
997 vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
998 const uint32_t *w, unsigned count)
999 {
1000 struct vtn_value *val = vtn_push_value(b, w[1], vtn_value_type_type);
1001
1002 val->type = rzalloc(b, struct vtn_type);
1003 val->type->id = w[1];
1004
1005 switch (opcode) {
1006 case SpvOpTypeVoid:
1007 val->type->base_type = vtn_base_type_void;
1008 val->type->type = glsl_void_type();
1009 break;
1010 case SpvOpTypeBool:
1011 val->type->base_type = vtn_base_type_scalar;
1012 val->type->type = glsl_bool_type();
1013 val->type->length = 1;
1014 break;
1015 case SpvOpTypeInt: {
1016 int bit_size = w[2];
1017 const bool signedness = w[3];
1018 val->type->base_type = vtn_base_type_scalar;
1019 switch (bit_size) {
1020 case 64:
1021 val->type->type = (signedness ? glsl_int64_t_type() : glsl_uint64_t_type());
1022 break;
1023 case 32:
1024 val->type->type = (signedness ? glsl_int_type() : glsl_uint_type());
1025 break;
1026 case 16:
1027 val->type->type = (signedness ? glsl_int16_t_type() : glsl_uint16_t_type());
1028 break;
1029 case 8:
1030 val->type->type = (signedness ? glsl_int8_t_type() : glsl_uint8_t_type());
1031 break;
1032 default:
1033 vtn_fail("Invalid int bit size");
1034 }
1035 val->type->length = 1;
1036 break;
1037 }
1038
1039 case SpvOpTypeFloat: {
1040 int bit_size = w[2];
1041 val->type->base_type = vtn_base_type_scalar;
1042 switch (bit_size) {
1043 case 16:
1044 val->type->type = glsl_float16_t_type();
1045 break;
1046 case 32:
1047 val->type->type = glsl_float_type();
1048 break;
1049 case 64:
1050 val->type->type = glsl_double_type();
1051 break;
1052 default:
1053 vtn_fail("Invalid float bit size");
1054 }
1055 val->type->length = 1;
1056 break;
1057 }
1058
1059 case SpvOpTypeVector: {
1060 struct vtn_type *base = vtn_value(b, w[2], vtn_value_type_type)->type;
1061 unsigned elems = w[3];
1062
1063 vtn_fail_if(base->base_type != vtn_base_type_scalar,
1064 "Base type for OpTypeVector must be a scalar");
1065 vtn_fail_if((elems < 2 || elems > 4) && (elems != 8) && (elems != 16),
1066 "Invalid component count for OpTypeVector");
1067
1068 val->type->base_type = vtn_base_type_vector;
1069 val->type->type = glsl_vector_type(glsl_get_base_type(base->type), elems);
1070 val->type->length = elems;
1071 val->type->stride = glsl_get_bit_size(base->type) / 8;
1072 val->type->array_element = base;
1073 break;
1074 }
1075
1076 case SpvOpTypeMatrix: {
1077 struct vtn_type *base = vtn_value(b, w[2], vtn_value_type_type)->type;
1078 unsigned columns = w[3];
1079
1080 vtn_fail_if(base->base_type != vtn_base_type_vector,
1081 "Base type for OpTypeMatrix must be a vector");
1082 vtn_fail_if(columns < 2 || columns > 4,
1083 "Invalid column count for OpTypeMatrix");
1084
1085 val->type->base_type = vtn_base_type_matrix;
1086 val->type->type = glsl_matrix_type(glsl_get_base_type(base->type),
1087 glsl_get_vector_elements(base->type),
1088 columns);
1089 vtn_fail_if(glsl_type_is_error(val->type->type),
1090 "Unsupported base type for OpTypeMatrix");
1091 assert(!glsl_type_is_error(val->type->type));
1092 val->type->length = columns;
1093 val->type->array_element = base;
1094 val->type->row_major = false;
1095 val->type->stride = 0;
1096 break;
1097 }
1098
1099 case SpvOpTypeRuntimeArray:
1100 case SpvOpTypeArray: {
1101 struct vtn_type *array_element =
1102 vtn_value(b, w[2], vtn_value_type_type)->type;
1103
1104 if (opcode == SpvOpTypeRuntimeArray) {
1105 /* A length of 0 is used to denote unsized arrays */
1106 val->type->length = 0;
1107 } else {
1108 val->type->length =
1109 vtn_value(b, w[3], vtn_value_type_constant)->constant->values[0].u32[0];
1110 }
1111
1112 val->type->base_type = vtn_base_type_array;
1113 val->type->type = glsl_array_type(array_element->type, val->type->length);
1114 val->type->array_element = array_element;
1115 val->type->stride = 0;
1116 break;
1117 }
1118
1119 case SpvOpTypeStruct: {
1120 unsigned num_fields = count - 2;
1121 val->type->base_type = vtn_base_type_struct;
1122 val->type->length = num_fields;
1123 val->type->members = ralloc_array(b, struct vtn_type *, num_fields);
1124 val->type->offsets = ralloc_array(b, unsigned, num_fields);
1125
1126 NIR_VLA(struct glsl_struct_field, fields, count);
1127 for (unsigned i = 0; i < num_fields; i++) {
1128 val->type->members[i] =
1129 vtn_value(b, w[i + 2], vtn_value_type_type)->type;
1130 fields[i] = (struct glsl_struct_field) {
1131 .type = val->type->members[i]->type,
1132 .name = ralloc_asprintf(b, "field%d", i),
1133 .location = -1,
1134 };
1135 }
1136
1137 struct member_decoration_ctx ctx = {
1138 .num_fields = num_fields,
1139 .fields = fields,
1140 .type = val->type
1141 };
1142
1143 vtn_foreach_decoration(b, val, struct_member_decoration_cb, &ctx);
1144 vtn_foreach_decoration(b, val, struct_member_matrix_stride_cb, &ctx);
1145
1146 const char *name = val->name ? val->name : "struct";
1147
1148 val->type->type = glsl_struct_type(fields, num_fields, name);
1149 break;
1150 }
1151
1152 case SpvOpTypeFunction: {
1153 val->type->base_type = vtn_base_type_function;
1154 val->type->type = NULL;
1155
1156 val->type->return_type = vtn_value(b, w[2], vtn_value_type_type)->type;
1157
1158 const unsigned num_params = count - 3;
1159 val->type->length = num_params;
1160 val->type->params = ralloc_array(b, struct vtn_type *, num_params);
1161 for (unsigned i = 0; i < count - 3; i++) {
1162 val->type->params[i] =
1163 vtn_value(b, w[i + 3], vtn_value_type_type)->type;
1164 }
1165 break;
1166 }
1167
1168 case SpvOpTypePointer: {
1169 SpvStorageClass storage_class = w[2];
1170 struct vtn_type *deref_type =
1171 vtn_value(b, w[3], vtn_value_type_type)->type;
1172
1173 val->type->base_type = vtn_base_type_pointer;
1174 val->type->storage_class = storage_class;
1175 val->type->deref = deref_type;
1176
1177 if (storage_class == SpvStorageClassUniform ||
1178 storage_class == SpvStorageClassStorageBuffer) {
1179 /* These can actually be stored to nir_variables and used as SSA
1180 * values so they need a real glsl_type.
1181 */
1182 val->type->type = glsl_vector_type(GLSL_TYPE_UINT, 2);
1183 }
1184
1185 if (storage_class == SpvStorageClassPushConstant) {
1186 /* These can actually be stored to nir_variables and used as SSA
1187 * values so they need a real glsl_type.
1188 */
1189 val->type->type = glsl_uint_type();
1190 }
1191
1192 if (storage_class == SpvStorageClassWorkgroup &&
1193 b->options->lower_workgroup_access_to_offsets) {
1194 uint32_t size, align;
1195 val->type->deref = vtn_type_layout_std430(b, val->type->deref,
1196 &size, &align);
1197 val->type->length = size;
1198 val->type->align = align;
1199 /* These can actually be stored to nir_variables and used as SSA
1200 * values so they need a real glsl_type.
1201 */
1202 val->type->type = glsl_uint_type();
1203 }
1204 break;
1205 }
1206
1207 case SpvOpTypeImage: {
1208 val->type->base_type = vtn_base_type_image;
1209
1210 const struct vtn_type *sampled_type =
1211 vtn_value(b, w[2], vtn_value_type_type)->type;
1212
1213 vtn_fail_if(sampled_type->base_type != vtn_base_type_scalar ||
1214 glsl_get_bit_size(sampled_type->type) != 32,
1215 "Sampled type of OpTypeImage must be a 32-bit scalar");
1216
1217 enum glsl_sampler_dim dim;
1218 switch ((SpvDim)w[3]) {
1219 case SpvDim1D: dim = GLSL_SAMPLER_DIM_1D; break;
1220 case SpvDim2D: dim = GLSL_SAMPLER_DIM_2D; break;
1221 case SpvDim3D: dim = GLSL_SAMPLER_DIM_3D; break;
1222 case SpvDimCube: dim = GLSL_SAMPLER_DIM_CUBE; break;
1223 case SpvDimRect: dim = GLSL_SAMPLER_DIM_RECT; break;
1224 case SpvDimBuffer: dim = GLSL_SAMPLER_DIM_BUF; break;
1225 case SpvDimSubpassData: dim = GLSL_SAMPLER_DIM_SUBPASS; break;
1226 default:
1227 vtn_fail("Invalid SPIR-V image dimensionality");
1228 }
1229
1230 bool is_shadow = w[4];
1231 bool is_array = w[5];
1232 bool multisampled = w[6];
1233 unsigned sampled = w[7];
1234 SpvImageFormat format = w[8];
1235
1236 if (count > 9)
1237 val->type->access_qualifier = w[9];
1238 else
1239 val->type->access_qualifier = SpvAccessQualifierReadWrite;
1240
1241 if (multisampled) {
1242 if (dim == GLSL_SAMPLER_DIM_2D)
1243 dim = GLSL_SAMPLER_DIM_MS;
1244 else if (dim == GLSL_SAMPLER_DIM_SUBPASS)
1245 dim = GLSL_SAMPLER_DIM_SUBPASS_MS;
1246 else
1247 vtn_fail("Unsupported multisampled image type");
1248 }
1249
1250 val->type->image_format = translate_image_format(b, format);
1251
1252 enum glsl_base_type sampled_base_type =
1253 glsl_get_base_type(sampled_type->type);
1254 if (sampled == 1) {
1255 val->type->sampled = true;
1256 val->type->type = glsl_sampler_type(dim, is_shadow, is_array,
1257 sampled_base_type);
1258 } else if (sampled == 2) {
1259 vtn_assert(!is_shadow);
1260 val->type->sampled = false;
1261 val->type->type = glsl_image_type(dim, is_array, sampled_base_type);
1262 } else {
1263 vtn_fail("We need to know if the image will be sampled");
1264 }
1265 break;
1266 }
1267
1268 case SpvOpTypeSampledImage:
1269 val->type->base_type = vtn_base_type_sampled_image;
1270 val->type->image = vtn_value(b, w[2], vtn_value_type_type)->type;
1271 val->type->type = val->type->image->type;
1272 break;
1273
1274 case SpvOpTypeSampler:
1275 /* The actual sampler type here doesn't really matter. It gets
1276 * thrown away the moment you combine it with an image. What really
1277 * matters is that it's a sampler type as opposed to an integer type
1278 * so the backend knows what to do.
1279 */
1280 val->type->base_type = vtn_base_type_sampler;
1281 val->type->type = glsl_bare_sampler_type();
1282 break;
1283
1284 case SpvOpTypeOpaque:
1285 case SpvOpTypeEvent:
1286 case SpvOpTypeDeviceEvent:
1287 case SpvOpTypeReserveId:
1288 case SpvOpTypeQueue:
1289 case SpvOpTypePipe:
1290 default:
1291 vtn_fail("Unhandled opcode");
1292 }
1293
1294 vtn_foreach_decoration(b, val, type_decoration_cb, NULL);
1295 }
1296
1297 static nir_constant *
1298 vtn_null_constant(struct vtn_builder *b, const struct glsl_type *type)
1299 {
1300 nir_constant *c = rzalloc(b, nir_constant);
1301
1302 /* For pointers and other typeless things, we have to return something but
1303 * it doesn't matter what.
1304 */
1305 if (!type)
1306 return c;
1307
1308 switch (glsl_get_base_type(type)) {
1309 case GLSL_TYPE_INT:
1310 case GLSL_TYPE_UINT:
1311 case GLSL_TYPE_INT16:
1312 case GLSL_TYPE_UINT16:
1313 case GLSL_TYPE_UINT8:
1314 case GLSL_TYPE_INT8:
1315 case GLSL_TYPE_INT64:
1316 case GLSL_TYPE_UINT64:
1317 case GLSL_TYPE_BOOL:
1318 case GLSL_TYPE_FLOAT:
1319 case GLSL_TYPE_FLOAT16:
1320 case GLSL_TYPE_DOUBLE:
1321 /* Nothing to do here. It's already initialized to zero */
1322 break;
1323
1324 case GLSL_TYPE_ARRAY:
1325 vtn_assert(glsl_get_length(type) > 0);
1326 c->num_elements = glsl_get_length(type);
1327 c->elements = ralloc_array(b, nir_constant *, c->num_elements);
1328
1329 c->elements[0] = vtn_null_constant(b, glsl_get_array_element(type));
1330 for (unsigned i = 1; i < c->num_elements; i++)
1331 c->elements[i] = c->elements[0];
1332 break;
1333
1334 case GLSL_TYPE_STRUCT:
1335 c->num_elements = glsl_get_length(type);
1336 c->elements = ralloc_array(b, nir_constant *, c->num_elements);
1337
1338 for (unsigned i = 0; i < c->num_elements; i++) {
1339 c->elements[i] = vtn_null_constant(b, glsl_get_struct_field(type, i));
1340 }
1341 break;
1342
1343 default:
1344 vtn_fail("Invalid type for null constant");
1345 }
1346
1347 return c;
1348 }
1349
1350 static void
1351 spec_constant_decoration_cb(struct vtn_builder *b, struct vtn_value *v,
1352 int member, const struct vtn_decoration *dec,
1353 void *data)
1354 {
1355 vtn_assert(member == -1);
1356 if (dec->decoration != SpvDecorationSpecId)
1357 return;
1358
1359 struct spec_constant_value *const_value = data;
1360
1361 for (unsigned i = 0; i < b->num_specializations; i++) {
1362 if (b->specializations[i].id == dec->literals[0]) {
1363 if (const_value->is_double)
1364 const_value->data64 = b->specializations[i].data64;
1365 else
1366 const_value->data32 = b->specializations[i].data32;
1367 return;
1368 }
1369 }
1370 }
1371
1372 static uint32_t
1373 get_specialization(struct vtn_builder *b, struct vtn_value *val,
1374 uint32_t const_value)
1375 {
1376 struct spec_constant_value data;
1377 data.is_double = false;
1378 data.data32 = const_value;
1379 vtn_foreach_decoration(b, val, spec_constant_decoration_cb, &data);
1380 return data.data32;
1381 }
1382
1383 static uint64_t
1384 get_specialization64(struct vtn_builder *b, struct vtn_value *val,
1385 uint64_t const_value)
1386 {
1387 struct spec_constant_value data;
1388 data.is_double = true;
1389 data.data64 = const_value;
1390 vtn_foreach_decoration(b, val, spec_constant_decoration_cb, &data);
1391 return data.data64;
1392 }
1393
1394 static void
1395 handle_workgroup_size_decoration_cb(struct vtn_builder *b,
1396 struct vtn_value *val,
1397 int member,
1398 const struct vtn_decoration *dec,
1399 void *data)
1400 {
1401 vtn_assert(member == -1);
1402 if (dec->decoration != SpvDecorationBuiltIn ||
1403 dec->literals[0] != SpvBuiltInWorkgroupSize)
1404 return;
1405
1406 vtn_assert(val->type->type == glsl_vector_type(GLSL_TYPE_UINT, 3));
1407
1408 b->shader->info.cs.local_size[0] = val->constant->values[0].u32[0];
1409 b->shader->info.cs.local_size[1] = val->constant->values[0].u32[1];
1410 b->shader->info.cs.local_size[2] = val->constant->values[0].u32[2];
1411 }
1412
1413 static void
1414 vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
1415 const uint32_t *w, unsigned count)
1416 {
1417 struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_constant);
1418 val->constant = rzalloc(b, nir_constant);
1419 switch (opcode) {
1420 case SpvOpConstantTrue:
1421 case SpvOpConstantFalse:
1422 case SpvOpSpecConstantTrue:
1423 case SpvOpSpecConstantFalse: {
1424 vtn_fail_if(val->type->type != glsl_bool_type(),
1425 "Result type of %s must be OpTypeBool",
1426 spirv_op_to_string(opcode));
1427
1428 uint32_t int_val = (opcode == SpvOpConstantTrue ||
1429 opcode == SpvOpSpecConstantTrue);
1430
1431 if (opcode == SpvOpSpecConstantTrue ||
1432 opcode == SpvOpSpecConstantFalse)
1433 int_val = get_specialization(b, val, int_val);
1434
1435 val->constant->values[0].u32[0] = int_val ? NIR_TRUE : NIR_FALSE;
1436 break;
1437 }
1438
1439 case SpvOpConstant: {
1440 vtn_fail_if(val->type->base_type != vtn_base_type_scalar,
1441 "Result type of %s must be a scalar",
1442 spirv_op_to_string(opcode));
1443 int bit_size = glsl_get_bit_size(val->type->type);
1444 switch (bit_size) {
1445 case 64:
1446 val->constant->values->u64[0] = vtn_u64_literal(&w[3]);
1447 break;
1448 case 32:
1449 val->constant->values->u32[0] = w[3];
1450 break;
1451 case 16:
1452 val->constant->values->u16[0] = w[3];
1453 break;
1454 case 8:
1455 val->constant->values->u8[0] = w[3];
1456 break;
1457 default:
1458 vtn_fail("Unsupported SpvOpConstant bit size");
1459 }
1460 break;
1461 }
1462
1463 case SpvOpSpecConstant: {
1464 vtn_fail_if(val->type->base_type != vtn_base_type_scalar,
1465 "Result type of %s must be a scalar",
1466 spirv_op_to_string(opcode));
1467 int bit_size = glsl_get_bit_size(val->type->type);
1468 switch (bit_size) {
1469 case 64:
1470 val->constant->values[0].u64[0] =
1471 get_specialization64(b, val, vtn_u64_literal(&w[3]));
1472 break;
1473 case 32:
1474 val->constant->values[0].u32[0] = get_specialization(b, val, w[3]);
1475 break;
1476 case 16:
1477 val->constant->values[0].u16[0] = get_specialization(b, val, w[3]);
1478 break;
1479 case 8:
1480 val->constant->values[0].u8[0] = get_specialization(b, val, w[3]);
1481 break;
1482 default:
1483 vtn_fail("Unsupported SpvOpSpecConstant bit size");
1484 }
1485 break;
1486 }
1487
1488 case SpvOpSpecConstantComposite:
1489 case SpvOpConstantComposite: {
1490 unsigned elem_count = count - 3;
1491 vtn_fail_if(elem_count != val->type->length,
1492 "%s has %u constituents, expected %u",
1493 spirv_op_to_string(opcode), elem_count, val->type->length);
1494
1495 nir_constant **elems = ralloc_array(b, nir_constant *, elem_count);
1496 for (unsigned i = 0; i < elem_count; i++)
1497 elems[i] = vtn_value(b, w[i + 3], vtn_value_type_constant)->constant;
1498
1499 switch (val->type->base_type) {
1500 case vtn_base_type_vector: {
1501 assert(glsl_type_is_vector(val->type->type));
1502 int bit_size = glsl_get_bit_size(val->type->type);
1503 for (unsigned i = 0; i < elem_count; i++) {
1504 switch (bit_size) {
1505 case 64:
1506 val->constant->values[0].u64[i] = elems[i]->values[0].u64[0];
1507 break;
1508 case 32:
1509 val->constant->values[0].u32[i] = elems[i]->values[0].u32[0];
1510 break;
1511 case 16:
1512 val->constant->values[0].u16[i] = elems[i]->values[0].u16[0];
1513 break;
1514 case 8:
1515 val->constant->values[0].u8[i] = elems[i]->values[0].u8[0];
1516 break;
1517 default:
1518 vtn_fail("Invalid SpvOpConstantComposite bit size");
1519 }
1520 }
1521 break;
1522 }
1523
1524 case vtn_base_type_matrix:
1525 assert(glsl_type_is_matrix(val->type->type));
1526 for (unsigned i = 0; i < elem_count; i++)
1527 val->constant->values[i] = elems[i]->values[0];
1528 break;
1529
1530 case vtn_base_type_struct:
1531 case vtn_base_type_array:
1532 ralloc_steal(val->constant, elems);
1533 val->constant->num_elements = elem_count;
1534 val->constant->elements = elems;
1535 break;
1536
1537 default:
1538 vtn_fail("Result type of %s must be a composite type",
1539 spirv_op_to_string(opcode));
1540 }
1541 break;
1542 }
1543
1544 case SpvOpSpecConstantOp: {
1545 SpvOp opcode = get_specialization(b, val, w[3]);
1546 switch (opcode) {
1547 case SpvOpVectorShuffle: {
1548 struct vtn_value *v0 = &b->values[w[4]];
1549 struct vtn_value *v1 = &b->values[w[5]];
1550
1551 vtn_assert(v0->value_type == vtn_value_type_constant ||
1552 v0->value_type == vtn_value_type_undef);
1553 vtn_assert(v1->value_type == vtn_value_type_constant ||
1554 v1->value_type == vtn_value_type_undef);
1555
1556 unsigned len0 = glsl_get_vector_elements(v0->type->type);
1557 unsigned len1 = glsl_get_vector_elements(v1->type->type);
1558
1559 vtn_assert(len0 + len1 < 16);
1560
1561 unsigned bit_size = glsl_get_bit_size(val->type->type);
1562 unsigned bit_size0 = glsl_get_bit_size(v0->type->type);
1563 unsigned bit_size1 = glsl_get_bit_size(v1->type->type);
1564
1565 vtn_assert(bit_size == bit_size0 && bit_size == bit_size1);
1566 (void)bit_size0; (void)bit_size1;
1567
1568 if (bit_size == 64) {
1569 uint64_t u64[8];
1570 if (v0->value_type == vtn_value_type_constant) {
1571 for (unsigned i = 0; i < len0; i++)
1572 u64[i] = v0->constant->values[0].u64[i];
1573 }
1574 if (v1->value_type == vtn_value_type_constant) {
1575 for (unsigned i = 0; i < len1; i++)
1576 u64[len0 + i] = v1->constant->values[0].u64[i];
1577 }
1578
1579 for (unsigned i = 0, j = 0; i < count - 6; i++, j++) {
1580 uint32_t comp = w[i + 6];
1581 /* If component is not used, set the value to a known constant
1582 * to detect if it is wrongly used.
1583 */
1584 if (comp == (uint32_t)-1)
1585 val->constant->values[0].u64[j] = 0xdeadbeefdeadbeef;
1586 else
1587 val->constant->values[0].u64[j] = u64[comp];
1588 }
1589 } else {
1590 /* This is for both 32-bit and 16-bit values */
1591 uint32_t u32[8];
1592 if (v0->value_type == vtn_value_type_constant) {
1593 for (unsigned i = 0; i < len0; i++)
1594 u32[i] = v0->constant->values[0].u32[i];
1595 }
1596 if (v1->value_type == vtn_value_type_constant) {
1597 for (unsigned i = 0; i < len1; i++)
1598 u32[len0 + i] = v1->constant->values[0].u32[i];
1599 }
1600
1601 for (unsigned i = 0, j = 0; i < count - 6; i++, j++) {
1602 uint32_t comp = w[i + 6];
1603 /* If component is not used, set the value to a known constant
1604 * to detect if it is wrongly used.
1605 */
1606 if (comp == (uint32_t)-1)
1607 val->constant->values[0].u32[j] = 0xdeadbeef;
1608 else
1609 val->constant->values[0].u32[j] = u32[comp];
1610 }
1611 }
1612 break;
1613 }
1614
1615 case SpvOpCompositeExtract:
1616 case SpvOpCompositeInsert: {
1617 struct vtn_value *comp;
1618 unsigned deref_start;
1619 struct nir_constant **c;
1620 if (opcode == SpvOpCompositeExtract) {
1621 comp = vtn_value(b, w[4], vtn_value_type_constant);
1622 deref_start = 5;
1623 c = &comp->constant;
1624 } else {
1625 comp = vtn_value(b, w[5], vtn_value_type_constant);
1626 deref_start = 6;
1627 val->constant = nir_constant_clone(comp->constant,
1628 (nir_variable *)b);
1629 c = &val->constant;
1630 }
1631
1632 int elem = -1;
1633 int col = 0;
1634 const struct vtn_type *type = comp->type;
1635 for (unsigned i = deref_start; i < count; i++) {
1636 vtn_fail_if(w[i] > type->length,
1637 "%uth index of %s is %u but the type has only "
1638 "%u elements", i - deref_start,
1639 spirv_op_to_string(opcode), w[i], type->length);
1640
1641 switch (type->base_type) {
1642 case vtn_base_type_vector:
1643 elem = w[i];
1644 type = type->array_element;
1645 break;
1646
1647 case vtn_base_type_matrix:
1648 assert(col == 0 && elem == -1);
1649 col = w[i];
1650 elem = 0;
1651 type = type->array_element;
1652 break;
1653
1654 case vtn_base_type_array:
1655 c = &(*c)->elements[w[i]];
1656 type = type->array_element;
1657 break;
1658
1659 case vtn_base_type_struct:
1660 c = &(*c)->elements[w[i]];
1661 type = type->members[w[i]];
1662 break;
1663
1664 default:
1665 vtn_fail("%s must only index into composite types",
1666 spirv_op_to_string(opcode));
1667 }
1668 }
1669
1670 if (opcode == SpvOpCompositeExtract) {
1671 if (elem == -1) {
1672 val->constant = *c;
1673 } else {
1674 unsigned num_components = type->length;
1675 unsigned bit_size = glsl_get_bit_size(type->type);
1676 for (unsigned i = 0; i < num_components; i++)
1677 switch(bit_size) {
1678 case 64:
1679 val->constant->values[0].u64[i] = (*c)->values[col].u64[elem + i];
1680 break;
1681 case 32:
1682 val->constant->values[0].u32[i] = (*c)->values[col].u32[elem + i];
1683 break;
1684 case 16:
1685 val->constant->values[0].u16[i] = (*c)->values[col].u16[elem + i];
1686 break;
1687 case 8:
1688 val->constant->values[0].u8[i] = (*c)->values[col].u8[elem + i];
1689 break;
1690 default:
1691 vtn_fail("Invalid SpvOpCompositeExtract bit size");
1692 }
1693 }
1694 } else {
1695 struct vtn_value *insert =
1696 vtn_value(b, w[4], vtn_value_type_constant);
1697 vtn_assert(insert->type == type);
1698 if (elem == -1) {
1699 *c = insert->constant;
1700 } else {
1701 unsigned num_components = type->length;
1702 unsigned bit_size = glsl_get_bit_size(type->type);
1703 for (unsigned i = 0; i < num_components; i++)
1704 switch (bit_size) {
1705 case 64:
1706 (*c)->values[col].u64[elem + i] = insert->constant->values[0].u64[i];
1707 break;
1708 case 32:
1709 (*c)->values[col].u32[elem + i] = insert->constant->values[0].u32[i];
1710 break;
1711 case 16:
1712 (*c)->values[col].u16[elem + i] = insert->constant->values[0].u16[i];
1713 break;
1714 case 8:
1715 (*c)->values[col].u8[elem + i] = insert->constant->values[0].u8[i];
1716 break;
1717 default:
1718 vtn_fail("Invalid SpvOpCompositeInsert bit size");
1719 }
1720 }
1721 }
1722 break;
1723 }
1724
1725 default: {
1726 bool swap;
1727 nir_alu_type dst_alu_type = nir_get_nir_type_for_glsl_type(val->type->type);
1728 nir_alu_type src_alu_type = dst_alu_type;
1729 unsigned num_components = glsl_get_vector_elements(val->type->type);
1730 unsigned bit_size;
1731
1732 vtn_assert(count <= 7);
1733
1734 switch (opcode) {
1735 case SpvOpSConvert:
1736 case SpvOpFConvert:
1737 /* We have a source in a conversion */
1738 src_alu_type =
1739 nir_get_nir_type_for_glsl_type(
1740 vtn_value(b, w[4], vtn_value_type_constant)->type->type);
1741 /* We use the bitsize of the conversion source to evaluate the opcode later */
1742 bit_size = glsl_get_bit_size(
1743 vtn_value(b, w[4], vtn_value_type_constant)->type->type);
1744 break;
1745 default:
1746 bit_size = glsl_get_bit_size(val->type->type);
1747 };
1748
1749 nir_op op = vtn_nir_alu_op_for_spirv_opcode(b, opcode, &swap,
1750 nir_alu_type_get_type_size(src_alu_type),
1751 nir_alu_type_get_type_size(dst_alu_type));
1752 nir_const_value src[4];
1753
1754 for (unsigned i = 0; i < count - 4; i++) {
1755 nir_constant *c =
1756 vtn_value(b, w[4 + i], vtn_value_type_constant)->constant;
1757
1758 unsigned j = swap ? 1 - i : i;
1759 src[j] = c->values[0];
1760 }
1761
1762 val->constant->values[0] =
1763 nir_eval_const_opcode(op, num_components, bit_size, src);
1764 break;
1765 } /* default */
1766 }
1767 break;
1768 }
1769
1770 case SpvOpConstantNull:
1771 val->constant = vtn_null_constant(b, val->type->type);
1772 break;
1773
1774 case SpvOpConstantSampler:
1775 vtn_fail("OpConstantSampler requires Kernel Capability");
1776 break;
1777
1778 default:
1779 vtn_fail("Unhandled opcode");
1780 }
1781
1782 /* Now that we have the value, update the workgroup size if needed */
1783 vtn_foreach_decoration(b, val, handle_workgroup_size_decoration_cb, NULL);
1784 }
1785
1786 static void
1787 vtn_handle_function_call(struct vtn_builder *b, SpvOp opcode,
1788 const uint32_t *w, unsigned count)
1789 {
1790 struct vtn_type *res_type = vtn_value(b, w[1], vtn_value_type_type)->type;
1791 struct vtn_function *vtn_callee =
1792 vtn_value(b, w[3], vtn_value_type_function)->func;
1793 struct nir_function *callee = vtn_callee->impl->function;
1794
1795 vtn_callee->referenced = true;
1796
1797 nir_call_instr *call = nir_call_instr_create(b->nb.shader, callee);
1798 for (unsigned i = 0; i < call->num_params; i++) {
1799 unsigned arg_id = w[4 + i];
1800 struct vtn_value *arg = vtn_untyped_value(b, arg_id);
1801 if (arg->value_type == vtn_value_type_pointer &&
1802 arg->pointer->ptr_type->type == NULL) {
1803 nir_deref_var *d = vtn_pointer_to_deref_var(b, arg->pointer);
1804 call->params[i] = nir_deref_var_clone(d, call);
1805 } else {
1806 struct vtn_ssa_value *arg_ssa = vtn_ssa_value(b, arg_id);
1807
1808 /* Make a temporary to store the argument in */
1809 nir_variable *tmp =
1810 nir_local_variable_create(b->nb.impl, arg_ssa->type, "arg_tmp");
1811 call->params[i] = nir_deref_var_create(call, tmp);
1812
1813 vtn_local_store(b, arg_ssa,
1814 nir_build_deref_for_chain(&b->nb, call->params[i]));
1815 }
1816 }
1817
1818 nir_variable *out_tmp = NULL;
1819 vtn_assert(res_type->type == callee->return_type);
1820 if (!glsl_type_is_void(callee->return_type)) {
1821 out_tmp = nir_local_variable_create(b->nb.impl, callee->return_type,
1822 "out_tmp");
1823 call->return_deref = nir_deref_var_create(call, out_tmp);
1824 }
1825
1826 nir_builder_instr_insert(&b->nb, &call->instr);
1827
1828 if (glsl_type_is_void(callee->return_type)) {
1829 vtn_push_value(b, w[2], vtn_value_type_undef);
1830 } else {
1831 nir_deref_instr *return_deref =
1832 nir_build_deref_for_chain(&b->nb, call->return_deref);
1833 vtn_push_ssa(b, w[2], res_type, vtn_local_load(b, return_deref));
1834 }
1835 }
1836
1837 struct vtn_ssa_value *
1838 vtn_create_ssa_value(struct vtn_builder *b, const struct glsl_type *type)
1839 {
1840 struct vtn_ssa_value *val = rzalloc(b, struct vtn_ssa_value);
1841 val->type = type;
1842
1843 if (!glsl_type_is_vector_or_scalar(type)) {
1844 unsigned elems = glsl_get_length(type);
1845 val->elems = ralloc_array(b, struct vtn_ssa_value *, elems);
1846 for (unsigned i = 0; i < elems; i++) {
1847 const struct glsl_type *child_type;
1848
1849 switch (glsl_get_base_type(type)) {
1850 case GLSL_TYPE_INT:
1851 case GLSL_TYPE_UINT:
1852 case GLSL_TYPE_INT16:
1853 case GLSL_TYPE_UINT16:
1854 case GLSL_TYPE_UINT8:
1855 case GLSL_TYPE_INT8:
1856 case GLSL_TYPE_INT64:
1857 case GLSL_TYPE_UINT64:
1858 case GLSL_TYPE_BOOL:
1859 case GLSL_TYPE_FLOAT:
1860 case GLSL_TYPE_FLOAT16:
1861 case GLSL_TYPE_DOUBLE:
1862 child_type = glsl_get_column_type(type);
1863 break;
1864 case GLSL_TYPE_ARRAY:
1865 child_type = glsl_get_array_element(type);
1866 break;
1867 case GLSL_TYPE_STRUCT:
1868 child_type = glsl_get_struct_field(type, i);
1869 break;
1870 default:
1871 vtn_fail("unkown base type");
1872 }
1873
1874 val->elems[i] = vtn_create_ssa_value(b, child_type);
1875 }
1876 }
1877
1878 return val;
1879 }
1880
1881 static nir_tex_src
1882 vtn_tex_src(struct vtn_builder *b, unsigned index, nir_tex_src_type type)
1883 {
1884 nir_tex_src src;
1885 src.src = nir_src_for_ssa(vtn_ssa_value(b, index)->def);
1886 src.src_type = type;
1887 return src;
1888 }
1889
1890 static void
1891 vtn_handle_texture(struct vtn_builder *b, SpvOp opcode,
1892 const uint32_t *w, unsigned count)
1893 {
1894 if (opcode == SpvOpSampledImage) {
1895 struct vtn_value *val =
1896 vtn_push_value(b, w[2], vtn_value_type_sampled_image);
1897 val->sampled_image = ralloc(b, struct vtn_sampled_image);
1898 val->sampled_image->type =
1899 vtn_value(b, w[1], vtn_value_type_type)->type;
1900 val->sampled_image->image =
1901 vtn_value(b, w[3], vtn_value_type_pointer)->pointer;
1902 val->sampled_image->sampler =
1903 vtn_value(b, w[4], vtn_value_type_pointer)->pointer;
1904 return;
1905 } else if (opcode == SpvOpImage) {
1906 struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_pointer);
1907 struct vtn_value *src_val = vtn_untyped_value(b, w[3]);
1908 if (src_val->value_type == vtn_value_type_sampled_image) {
1909 val->pointer = src_val->sampled_image->image;
1910 } else {
1911 vtn_assert(src_val->value_type == vtn_value_type_pointer);
1912 val->pointer = src_val->pointer;
1913 }
1914 return;
1915 }
1916
1917 struct vtn_type *ret_type = vtn_value(b, w[1], vtn_value_type_type)->type;
1918 struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_ssa);
1919
1920 struct vtn_sampled_image sampled;
1921 struct vtn_value *sampled_val = vtn_untyped_value(b, w[3]);
1922 if (sampled_val->value_type == vtn_value_type_sampled_image) {
1923 sampled = *sampled_val->sampled_image;
1924 } else {
1925 vtn_assert(sampled_val->value_type == vtn_value_type_pointer);
1926 sampled.type = sampled_val->pointer->type;
1927 sampled.image = NULL;
1928 sampled.sampler = sampled_val->pointer;
1929 }
1930
1931 const struct glsl_type *image_type = sampled.type->type;
1932 const enum glsl_sampler_dim sampler_dim = glsl_get_sampler_dim(image_type);
1933 const bool is_array = glsl_sampler_type_is_array(image_type);
1934
1935 /* Figure out the base texture operation */
1936 nir_texop texop;
1937 switch (opcode) {
1938 case SpvOpImageSampleImplicitLod:
1939 case SpvOpImageSampleDrefImplicitLod:
1940 case SpvOpImageSampleProjImplicitLod:
1941 case SpvOpImageSampleProjDrefImplicitLod:
1942 texop = nir_texop_tex;
1943 break;
1944
1945 case SpvOpImageSampleExplicitLod:
1946 case SpvOpImageSampleDrefExplicitLod:
1947 case SpvOpImageSampleProjExplicitLod:
1948 case SpvOpImageSampleProjDrefExplicitLod:
1949 texop = nir_texop_txl;
1950 break;
1951
1952 case SpvOpImageFetch:
1953 if (glsl_get_sampler_dim(image_type) == GLSL_SAMPLER_DIM_MS) {
1954 texop = nir_texop_txf_ms;
1955 } else {
1956 texop = nir_texop_txf;
1957 }
1958 break;
1959
1960 case SpvOpImageGather:
1961 case SpvOpImageDrefGather:
1962 texop = nir_texop_tg4;
1963 break;
1964
1965 case SpvOpImageQuerySizeLod:
1966 case SpvOpImageQuerySize:
1967 texop = nir_texop_txs;
1968 break;
1969
1970 case SpvOpImageQueryLod:
1971 texop = nir_texop_lod;
1972 break;
1973
1974 case SpvOpImageQueryLevels:
1975 texop = nir_texop_query_levels;
1976 break;
1977
1978 case SpvOpImageQuerySamples:
1979 texop = nir_texop_texture_samples;
1980 break;
1981
1982 default:
1983 vtn_fail("Unhandled opcode");
1984 }
1985
1986 nir_tex_src srcs[10]; /* 10 should be enough */
1987 nir_tex_src *p = srcs;
1988
1989 nir_deref_instr *sampler = vtn_pointer_to_deref(b, sampled.sampler);
1990 nir_deref_instr *texture =
1991 sampled.image ? vtn_pointer_to_deref(b, sampled.image) : sampler;
1992
1993 p->src = nir_src_for_ssa(&texture->dest.ssa);
1994 p->src_type = nir_tex_src_texture_deref;
1995 p++;
1996
1997 switch (texop) {
1998 case nir_texop_tex:
1999 case nir_texop_txb:
2000 case nir_texop_txl:
2001 case nir_texop_txd:
2002 case nir_texop_tg4:
2003 /* These operations require a sampler */
2004 p->src = nir_src_for_ssa(&sampler->dest.ssa);
2005 p->src_type = nir_tex_src_sampler_deref;
2006 p++;
2007 break;
2008 case nir_texop_txf:
2009 case nir_texop_txf_ms:
2010 case nir_texop_txs:
2011 case nir_texop_lod:
2012 case nir_texop_query_levels:
2013 case nir_texop_texture_samples:
2014 case nir_texop_samples_identical:
2015 /* These don't */
2016 break;
2017 case nir_texop_txf_ms_mcs:
2018 vtn_fail("unexpected nir_texop_txf_ms_mcs");
2019 }
2020
2021 unsigned idx = 4;
2022
2023 struct nir_ssa_def *coord;
2024 unsigned coord_components;
2025 switch (opcode) {
2026 case SpvOpImageSampleImplicitLod:
2027 case SpvOpImageSampleExplicitLod:
2028 case SpvOpImageSampleDrefImplicitLod:
2029 case SpvOpImageSampleDrefExplicitLod:
2030 case SpvOpImageSampleProjImplicitLod:
2031 case SpvOpImageSampleProjExplicitLod:
2032 case SpvOpImageSampleProjDrefImplicitLod:
2033 case SpvOpImageSampleProjDrefExplicitLod:
2034 case SpvOpImageFetch:
2035 case SpvOpImageGather:
2036 case SpvOpImageDrefGather:
2037 case SpvOpImageQueryLod: {
2038 /* All these types have the coordinate as their first real argument */
2039 switch (sampler_dim) {
2040 case GLSL_SAMPLER_DIM_1D:
2041 case GLSL_SAMPLER_DIM_BUF:
2042 coord_components = 1;
2043 break;
2044 case GLSL_SAMPLER_DIM_2D:
2045 case GLSL_SAMPLER_DIM_RECT:
2046 case GLSL_SAMPLER_DIM_MS:
2047 coord_components = 2;
2048 break;
2049 case GLSL_SAMPLER_DIM_3D:
2050 case GLSL_SAMPLER_DIM_CUBE:
2051 coord_components = 3;
2052 break;
2053 default:
2054 vtn_fail("Invalid sampler type");
2055 }
2056
2057 if (is_array && texop != nir_texop_lod)
2058 coord_components++;
2059
2060 coord = vtn_ssa_value(b, w[idx++])->def;
2061 p->src = nir_src_for_ssa(nir_channels(&b->nb, coord,
2062 (1 << coord_components) - 1));
2063 p->src_type = nir_tex_src_coord;
2064 p++;
2065 break;
2066 }
2067
2068 default:
2069 coord = NULL;
2070 coord_components = 0;
2071 break;
2072 }
2073
2074 switch (opcode) {
2075 case SpvOpImageSampleProjImplicitLod:
2076 case SpvOpImageSampleProjExplicitLod:
2077 case SpvOpImageSampleProjDrefImplicitLod:
2078 case SpvOpImageSampleProjDrefExplicitLod:
2079 /* These have the projector as the last coordinate component */
2080 p->src = nir_src_for_ssa(nir_channel(&b->nb, coord, coord_components));
2081 p->src_type = nir_tex_src_projector;
2082 p++;
2083 break;
2084
2085 default:
2086 break;
2087 }
2088
2089 bool is_shadow = false;
2090 unsigned gather_component = 0;
2091 switch (opcode) {
2092 case SpvOpImageSampleDrefImplicitLod:
2093 case SpvOpImageSampleDrefExplicitLod:
2094 case SpvOpImageSampleProjDrefImplicitLod:
2095 case SpvOpImageSampleProjDrefExplicitLod:
2096 case SpvOpImageDrefGather:
2097 /* These all have an explicit depth value as their next source */
2098 is_shadow = true;
2099 (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_comparator);
2100 break;
2101
2102 case SpvOpImageGather:
2103 /* This has a component as its next source */
2104 gather_component =
2105 vtn_value(b, w[idx++], vtn_value_type_constant)->constant->values[0].u32[0];
2106 break;
2107
2108 default:
2109 break;
2110 }
2111
2112 /* For OpImageQuerySizeLod, we always have an LOD */
2113 if (opcode == SpvOpImageQuerySizeLod)
2114 (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_lod);
2115
2116 /* Now we need to handle some number of optional arguments */
2117 const struct vtn_ssa_value *gather_offsets = NULL;
2118 if (idx < count) {
2119 uint32_t operands = w[idx++];
2120
2121 if (operands & SpvImageOperandsBiasMask) {
2122 vtn_assert(texop == nir_texop_tex);
2123 texop = nir_texop_txb;
2124 (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_bias);
2125 }
2126
2127 if (operands & SpvImageOperandsLodMask) {
2128 vtn_assert(texop == nir_texop_txl || texop == nir_texop_txf ||
2129 texop == nir_texop_txs);
2130 (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_lod);
2131 }
2132
2133 if (operands & SpvImageOperandsGradMask) {
2134 vtn_assert(texop == nir_texop_txl);
2135 texop = nir_texop_txd;
2136 (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_ddx);
2137 (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_ddy);
2138 }
2139
2140 if (operands & SpvImageOperandsOffsetMask ||
2141 operands & SpvImageOperandsConstOffsetMask)
2142 (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_offset);
2143
2144 if (operands & SpvImageOperandsConstOffsetsMask) {
2145 nir_tex_src none = {0};
2146 gather_offsets = vtn_ssa_value(b, w[idx++]);
2147 (*p++) = none;
2148 }
2149
2150 if (operands & SpvImageOperandsSampleMask) {
2151 vtn_assert(texop == nir_texop_txf_ms);
2152 texop = nir_texop_txf_ms;
2153 (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_ms_index);
2154 }
2155 }
2156 /* We should have now consumed exactly all of the arguments */
2157 vtn_assert(idx == count);
2158
2159 nir_tex_instr *instr = nir_tex_instr_create(b->shader, p - srcs);
2160 instr->op = texop;
2161
2162 memcpy(instr->src, srcs, instr->num_srcs * sizeof(*instr->src));
2163
2164 instr->coord_components = coord_components;
2165 instr->sampler_dim = sampler_dim;
2166 instr->is_array = is_array;
2167 instr->is_shadow = is_shadow;
2168 instr->is_new_style_shadow =
2169 is_shadow && glsl_get_components(ret_type->type) == 1;
2170 instr->component = gather_component;
2171
2172 switch (glsl_get_sampler_result_type(image_type)) {
2173 case GLSL_TYPE_FLOAT: instr->dest_type = nir_type_float; break;
2174 case GLSL_TYPE_INT: instr->dest_type = nir_type_int; break;
2175 case GLSL_TYPE_UINT: instr->dest_type = nir_type_uint; break;
2176 case GLSL_TYPE_BOOL: instr->dest_type = nir_type_bool; break;
2177 default:
2178 vtn_fail("Invalid base type for sampler result");
2179 }
2180
2181 nir_ssa_dest_init(&instr->instr, &instr->dest,
2182 nir_tex_instr_dest_size(instr), 32, NULL);
2183
2184 vtn_assert(glsl_get_vector_elements(ret_type->type) ==
2185 nir_tex_instr_dest_size(instr));
2186
2187 nir_ssa_def *def;
2188 nir_instr *instruction;
2189 if (gather_offsets) {
2190 vtn_assert(glsl_get_base_type(gather_offsets->type) == GLSL_TYPE_ARRAY);
2191 vtn_assert(glsl_get_length(gather_offsets->type) == 4);
2192 nir_tex_instr *instrs[4] = {instr, NULL, NULL, NULL};
2193
2194 /* Copy the current instruction 4x */
2195 for (uint32_t i = 1; i < 4; i++) {
2196 instrs[i] = nir_tex_instr_create(b->shader, instr->num_srcs);
2197 instrs[i]->op = instr->op;
2198 instrs[i]->coord_components = instr->coord_components;
2199 instrs[i]->sampler_dim = instr->sampler_dim;
2200 instrs[i]->is_array = instr->is_array;
2201 instrs[i]->is_shadow = instr->is_shadow;
2202 instrs[i]->is_new_style_shadow = instr->is_new_style_shadow;
2203 instrs[i]->component = instr->component;
2204 instrs[i]->dest_type = instr->dest_type;
2205
2206 memcpy(instrs[i]->src, srcs, instr->num_srcs * sizeof(*instr->src));
2207
2208 nir_ssa_dest_init(&instrs[i]->instr, &instrs[i]->dest,
2209 nir_tex_instr_dest_size(instr), 32, NULL);
2210 }
2211
2212 /* Fill in the last argument with the offset from the passed in offsets
2213 * and insert the instruction into the stream.
2214 */
2215 for (uint32_t i = 0; i < 4; i++) {
2216 nir_tex_src src;
2217 src.src = nir_src_for_ssa(gather_offsets->elems[i]->def);
2218 src.src_type = nir_tex_src_offset;
2219 instrs[i]->src[instrs[i]->num_srcs - 1] = src;
2220 nir_builder_instr_insert(&b->nb, &instrs[i]->instr);
2221 }
2222
2223 /* Combine the results of the 4 instructions by taking their .w
2224 * components
2225 */
2226 nir_alu_instr *vec4 = nir_alu_instr_create(b->shader, nir_op_vec4);
2227 nir_ssa_dest_init(&vec4->instr, &vec4->dest.dest, 4, 32, NULL);
2228 vec4->dest.write_mask = 0xf;
2229 for (uint32_t i = 0; i < 4; i++) {
2230 vec4->src[i].src = nir_src_for_ssa(&instrs[i]->dest.ssa);
2231 vec4->src[i].swizzle[0] = 3;
2232 }
2233 def = &vec4->dest.dest.ssa;
2234 instruction = &vec4->instr;
2235 } else {
2236 def = &instr->dest.ssa;
2237 instruction = &instr->instr;
2238 }
2239
2240 val->ssa = vtn_create_ssa_value(b, ret_type->type);
2241 val->ssa->def = def;
2242
2243 nir_builder_instr_insert(&b->nb, instruction);
2244 }
2245
2246 static void
2247 fill_common_atomic_sources(struct vtn_builder *b, SpvOp opcode,
2248 const uint32_t *w, nir_src *src)
2249 {
2250 switch (opcode) {
2251 case SpvOpAtomicIIncrement:
2252 src[0] = nir_src_for_ssa(nir_imm_int(&b->nb, 1));
2253 break;
2254
2255 case SpvOpAtomicIDecrement:
2256 src[0] = nir_src_for_ssa(nir_imm_int(&b->nb, -1));
2257 break;
2258
2259 case SpvOpAtomicISub:
2260 src[0] =
2261 nir_src_for_ssa(nir_ineg(&b->nb, vtn_ssa_value(b, w[6])->def));
2262 break;
2263
2264 case SpvOpAtomicCompareExchange:
2265 src[0] = nir_src_for_ssa(vtn_ssa_value(b, w[8])->def);
2266 src[1] = nir_src_for_ssa(vtn_ssa_value(b, w[7])->def);
2267 break;
2268
2269 case SpvOpAtomicExchange:
2270 case SpvOpAtomicIAdd:
2271 case SpvOpAtomicSMin:
2272 case SpvOpAtomicUMin:
2273 case SpvOpAtomicSMax:
2274 case SpvOpAtomicUMax:
2275 case SpvOpAtomicAnd:
2276 case SpvOpAtomicOr:
2277 case SpvOpAtomicXor:
2278 src[0] = nir_src_for_ssa(vtn_ssa_value(b, w[6])->def);
2279 break;
2280
2281 default:
2282 vtn_fail("Invalid SPIR-V atomic");
2283 }
2284 }
2285
2286 static nir_ssa_def *
2287 get_image_coord(struct vtn_builder *b, uint32_t value)
2288 {
2289 struct vtn_ssa_value *coord = vtn_ssa_value(b, value);
2290
2291 /* The image_load_store intrinsics assume a 4-dim coordinate */
2292 unsigned dim = glsl_get_vector_elements(coord->type);
2293 unsigned swizzle[4];
2294 for (unsigned i = 0; i < 4; i++)
2295 swizzle[i] = MIN2(i, dim - 1);
2296
2297 return nir_swizzle(&b->nb, coord->def, swizzle, 4, false);
2298 }
2299
2300 static void
2301 vtn_handle_image(struct vtn_builder *b, SpvOp opcode,
2302 const uint32_t *w, unsigned count)
2303 {
2304 /* Just get this one out of the way */
2305 if (opcode == SpvOpImageTexelPointer) {
2306 struct vtn_value *val =
2307 vtn_push_value(b, w[2], vtn_value_type_image_pointer);
2308 val->image = ralloc(b, struct vtn_image_pointer);
2309
2310 val->image->image = vtn_value(b, w[3], vtn_value_type_pointer)->pointer;
2311 val->image->coord = get_image_coord(b, w[4]);
2312 val->image->sample = vtn_ssa_value(b, w[5])->def;
2313 return;
2314 }
2315
2316 struct vtn_image_pointer image;
2317
2318 switch (opcode) {
2319 case SpvOpAtomicExchange:
2320 case SpvOpAtomicCompareExchange:
2321 case SpvOpAtomicCompareExchangeWeak:
2322 case SpvOpAtomicIIncrement:
2323 case SpvOpAtomicIDecrement:
2324 case SpvOpAtomicIAdd:
2325 case SpvOpAtomicISub:
2326 case SpvOpAtomicLoad:
2327 case SpvOpAtomicSMin:
2328 case SpvOpAtomicUMin:
2329 case SpvOpAtomicSMax:
2330 case SpvOpAtomicUMax:
2331 case SpvOpAtomicAnd:
2332 case SpvOpAtomicOr:
2333 case SpvOpAtomicXor:
2334 image = *vtn_value(b, w[3], vtn_value_type_image_pointer)->image;
2335 break;
2336
2337 case SpvOpAtomicStore:
2338 image = *vtn_value(b, w[1], vtn_value_type_image_pointer)->image;
2339 break;
2340
2341 case SpvOpImageQuerySize:
2342 image.image = vtn_value(b, w[3], vtn_value_type_pointer)->pointer;
2343 image.coord = NULL;
2344 image.sample = NULL;
2345 break;
2346
2347 case SpvOpImageRead:
2348 image.image = vtn_value(b, w[3], vtn_value_type_pointer)->pointer;
2349 image.coord = get_image_coord(b, w[4]);
2350
2351 if (count > 5 && (w[5] & SpvImageOperandsSampleMask)) {
2352 vtn_assert(w[5] == SpvImageOperandsSampleMask);
2353 image.sample = vtn_ssa_value(b, w[6])->def;
2354 } else {
2355 image.sample = nir_ssa_undef(&b->nb, 1, 32);
2356 }
2357 break;
2358
2359 case SpvOpImageWrite:
2360 image.image = vtn_value(b, w[1], vtn_value_type_pointer)->pointer;
2361 image.coord = get_image_coord(b, w[2]);
2362
2363 /* texel = w[3] */
2364
2365 if (count > 4 && (w[4] & SpvImageOperandsSampleMask)) {
2366 vtn_assert(w[4] == SpvImageOperandsSampleMask);
2367 image.sample = vtn_ssa_value(b, w[5])->def;
2368 } else {
2369 image.sample = nir_ssa_undef(&b->nb, 1, 32);
2370 }
2371 break;
2372
2373 default:
2374 vtn_fail("Invalid image opcode");
2375 }
2376
2377 nir_intrinsic_op op;
2378 switch (opcode) {
2379 #define OP(S, N) case SpvOp##S: op = nir_intrinsic_image_deref_##N; break;
2380 OP(ImageQuerySize, size)
2381 OP(ImageRead, load)
2382 OP(ImageWrite, store)
2383 OP(AtomicLoad, load)
2384 OP(AtomicStore, store)
2385 OP(AtomicExchange, atomic_exchange)
2386 OP(AtomicCompareExchange, atomic_comp_swap)
2387 OP(AtomicIIncrement, atomic_add)
2388 OP(AtomicIDecrement, atomic_add)
2389 OP(AtomicIAdd, atomic_add)
2390 OP(AtomicISub, atomic_add)
2391 OP(AtomicSMin, atomic_min)
2392 OP(AtomicUMin, atomic_min)
2393 OP(AtomicSMax, atomic_max)
2394 OP(AtomicUMax, atomic_max)
2395 OP(AtomicAnd, atomic_and)
2396 OP(AtomicOr, atomic_or)
2397 OP(AtomicXor, atomic_xor)
2398 #undef OP
2399 default:
2400 vtn_fail("Invalid image opcode");
2401 }
2402
2403 nir_intrinsic_instr *intrin = nir_intrinsic_instr_create(b->shader, op);
2404
2405 nir_deref_instr *image_deref = vtn_pointer_to_deref(b, image.image);
2406 intrin->src[0] = nir_src_for_ssa(&image_deref->dest.ssa);
2407
2408 /* ImageQuerySize doesn't take any extra parameters */
2409 if (opcode != SpvOpImageQuerySize) {
2410 /* The image coordinate is always 4 components but we may not have that
2411 * many. Swizzle to compensate.
2412 */
2413 unsigned swiz[4];
2414 for (unsigned i = 0; i < 4; i++)
2415 swiz[i] = i < image.coord->num_components ? i : 0;
2416 intrin->src[1] = nir_src_for_ssa(nir_swizzle(&b->nb, image.coord,
2417 swiz, 4, false));
2418 intrin->src[2] = nir_src_for_ssa(image.sample);
2419 }
2420
2421 switch (opcode) {
2422 case SpvOpAtomicLoad:
2423 case SpvOpImageQuerySize:
2424 case SpvOpImageRead:
2425 break;
2426 case SpvOpAtomicStore:
2427 intrin->src[3] = nir_src_for_ssa(vtn_ssa_value(b, w[4])->def);
2428 break;
2429 case SpvOpImageWrite:
2430 intrin->src[3] = nir_src_for_ssa(vtn_ssa_value(b, w[3])->def);
2431 break;
2432
2433 case SpvOpAtomicCompareExchange:
2434 case SpvOpAtomicIIncrement:
2435 case SpvOpAtomicIDecrement:
2436 case SpvOpAtomicExchange:
2437 case SpvOpAtomicIAdd:
2438 case SpvOpAtomicISub:
2439 case SpvOpAtomicSMin:
2440 case SpvOpAtomicUMin:
2441 case SpvOpAtomicSMax:
2442 case SpvOpAtomicUMax:
2443 case SpvOpAtomicAnd:
2444 case SpvOpAtomicOr:
2445 case SpvOpAtomicXor:
2446 fill_common_atomic_sources(b, opcode, w, &intrin->src[3]);
2447 break;
2448
2449 default:
2450 vtn_fail("Invalid image opcode");
2451 }
2452
2453 if (opcode != SpvOpImageWrite) {
2454 struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_ssa);
2455 struct vtn_type *type = vtn_value(b, w[1], vtn_value_type_type)->type;
2456
2457 unsigned dest_components = nir_intrinsic_dest_components(intrin);
2458 if (intrin->intrinsic == nir_intrinsic_image_deref_size) {
2459 dest_components = intrin->num_components =
2460 glsl_get_vector_elements(type->type);
2461 }
2462
2463 nir_ssa_dest_init(&intrin->instr, &intrin->dest,
2464 dest_components, 32, NULL);
2465
2466 nir_builder_instr_insert(&b->nb, &intrin->instr);
2467
2468 val->ssa = vtn_create_ssa_value(b, type->type);
2469 val->ssa->def = &intrin->dest.ssa;
2470 } else {
2471 nir_builder_instr_insert(&b->nb, &intrin->instr);
2472 }
2473 }
2474
2475 static nir_intrinsic_op
2476 get_ssbo_nir_atomic_op(struct vtn_builder *b, SpvOp opcode)
2477 {
2478 switch (opcode) {
2479 case SpvOpAtomicLoad: return nir_intrinsic_load_ssbo;
2480 case SpvOpAtomicStore: return nir_intrinsic_store_ssbo;
2481 #define OP(S, N) case SpvOp##S: return nir_intrinsic_ssbo_##N;
2482 OP(AtomicExchange, atomic_exchange)
2483 OP(AtomicCompareExchange, atomic_comp_swap)
2484 OP(AtomicIIncrement, atomic_add)
2485 OP(AtomicIDecrement, atomic_add)
2486 OP(AtomicIAdd, atomic_add)
2487 OP(AtomicISub, atomic_add)
2488 OP(AtomicSMin, atomic_imin)
2489 OP(AtomicUMin, atomic_umin)
2490 OP(AtomicSMax, atomic_imax)
2491 OP(AtomicUMax, atomic_umax)
2492 OP(AtomicAnd, atomic_and)
2493 OP(AtomicOr, atomic_or)
2494 OP(AtomicXor, atomic_xor)
2495 #undef OP
2496 default:
2497 vtn_fail("Invalid SSBO atomic");
2498 }
2499 }
2500
2501 static nir_intrinsic_op
2502 get_shared_nir_atomic_op(struct vtn_builder *b, SpvOp opcode)
2503 {
2504 switch (opcode) {
2505 case SpvOpAtomicLoad: return nir_intrinsic_load_shared;
2506 case SpvOpAtomicStore: return nir_intrinsic_store_shared;
2507 #define OP(S, N) case SpvOp##S: return nir_intrinsic_shared_##N;
2508 OP(AtomicExchange, atomic_exchange)
2509 OP(AtomicCompareExchange, atomic_comp_swap)
2510 OP(AtomicIIncrement, atomic_add)
2511 OP(AtomicIDecrement, atomic_add)
2512 OP(AtomicIAdd, atomic_add)
2513 OP(AtomicISub, atomic_add)
2514 OP(AtomicSMin, atomic_imin)
2515 OP(AtomicUMin, atomic_umin)
2516 OP(AtomicSMax, atomic_imax)
2517 OP(AtomicUMax, atomic_umax)
2518 OP(AtomicAnd, atomic_and)
2519 OP(AtomicOr, atomic_or)
2520 OP(AtomicXor, atomic_xor)
2521 #undef OP
2522 default:
2523 vtn_fail("Invalid shared atomic");
2524 }
2525 }
2526
2527 static nir_intrinsic_op
2528 get_deref_nir_atomic_op(struct vtn_builder *b, SpvOp opcode)
2529 {
2530 switch (opcode) {
2531 case SpvOpAtomicLoad: return nir_intrinsic_load_deref;
2532 case SpvOpAtomicStore: return nir_intrinsic_store_deref;
2533 #define OP(S, N) case SpvOp##S: return nir_intrinsic_deref_##N;
2534 OP(AtomicExchange, atomic_exchange)
2535 OP(AtomicCompareExchange, atomic_comp_swap)
2536 OP(AtomicIIncrement, atomic_add)
2537 OP(AtomicIDecrement, atomic_add)
2538 OP(AtomicIAdd, atomic_add)
2539 OP(AtomicISub, atomic_add)
2540 OP(AtomicSMin, atomic_imin)
2541 OP(AtomicUMin, atomic_umin)
2542 OP(AtomicSMax, atomic_imax)
2543 OP(AtomicUMax, atomic_umax)
2544 OP(AtomicAnd, atomic_and)
2545 OP(AtomicOr, atomic_or)
2546 OP(AtomicXor, atomic_xor)
2547 #undef OP
2548 default:
2549 vtn_fail("Invalid shared atomic");
2550 }
2551 }
2552
2553 static void
2554 vtn_handle_ssbo_or_shared_atomic(struct vtn_builder *b, SpvOp opcode,
2555 const uint32_t *w, unsigned count)
2556 {
2557 struct vtn_pointer *ptr;
2558 nir_intrinsic_instr *atomic;
2559
2560 switch (opcode) {
2561 case SpvOpAtomicLoad:
2562 case SpvOpAtomicExchange:
2563 case SpvOpAtomicCompareExchange:
2564 case SpvOpAtomicCompareExchangeWeak:
2565 case SpvOpAtomicIIncrement:
2566 case SpvOpAtomicIDecrement:
2567 case SpvOpAtomicIAdd:
2568 case SpvOpAtomicISub:
2569 case SpvOpAtomicSMin:
2570 case SpvOpAtomicUMin:
2571 case SpvOpAtomicSMax:
2572 case SpvOpAtomicUMax:
2573 case SpvOpAtomicAnd:
2574 case SpvOpAtomicOr:
2575 case SpvOpAtomicXor:
2576 ptr = vtn_value(b, w[3], vtn_value_type_pointer)->pointer;
2577 break;
2578
2579 case SpvOpAtomicStore:
2580 ptr = vtn_value(b, w[1], vtn_value_type_pointer)->pointer;
2581 break;
2582
2583 default:
2584 vtn_fail("Invalid SPIR-V atomic");
2585 }
2586
2587 /*
2588 SpvScope scope = w[4];
2589 SpvMemorySemanticsMask semantics = w[5];
2590 */
2591
2592 if (ptr->mode == vtn_variable_mode_workgroup &&
2593 !b->options->lower_workgroup_access_to_offsets) {
2594 nir_deref_instr *deref = vtn_pointer_to_deref(b, ptr);
2595 const struct glsl_type *deref_type = deref->type;
2596 nir_intrinsic_op op = get_deref_nir_atomic_op(b, opcode);
2597 atomic = nir_intrinsic_instr_create(b->nb.shader, op);
2598 atomic->src[0] = nir_src_for_ssa(&deref->dest.ssa);
2599
2600 switch (opcode) {
2601 case SpvOpAtomicLoad:
2602 atomic->num_components = glsl_get_vector_elements(deref_type);
2603 break;
2604
2605 case SpvOpAtomicStore:
2606 atomic->num_components = glsl_get_vector_elements(deref_type);
2607 nir_intrinsic_set_write_mask(atomic, (1 << atomic->num_components) - 1);
2608 atomic->src[1] = nir_src_for_ssa(vtn_ssa_value(b, w[4])->def);
2609 break;
2610
2611 case SpvOpAtomicExchange:
2612 case SpvOpAtomicCompareExchange:
2613 case SpvOpAtomicCompareExchangeWeak:
2614 case SpvOpAtomicIIncrement:
2615 case SpvOpAtomicIDecrement:
2616 case SpvOpAtomicIAdd:
2617 case SpvOpAtomicISub:
2618 case SpvOpAtomicSMin:
2619 case SpvOpAtomicUMin:
2620 case SpvOpAtomicSMax:
2621 case SpvOpAtomicUMax:
2622 case SpvOpAtomicAnd:
2623 case SpvOpAtomicOr:
2624 case SpvOpAtomicXor:
2625 fill_common_atomic_sources(b, opcode, w, &atomic->src[1]);
2626 break;
2627
2628 default:
2629 vtn_fail("Invalid SPIR-V atomic");
2630
2631 }
2632 } else {
2633 nir_ssa_def *offset, *index;
2634 offset = vtn_pointer_to_offset(b, ptr, &index, NULL);
2635
2636 nir_intrinsic_op op;
2637 if (ptr->mode == vtn_variable_mode_ssbo) {
2638 op = get_ssbo_nir_atomic_op(b, opcode);
2639 } else {
2640 vtn_assert(ptr->mode == vtn_variable_mode_workgroup &&
2641 b->options->lower_workgroup_access_to_offsets);
2642 op = get_shared_nir_atomic_op(b, opcode);
2643 }
2644
2645 atomic = nir_intrinsic_instr_create(b->nb.shader, op);
2646
2647 int src = 0;
2648 switch (opcode) {
2649 case SpvOpAtomicLoad:
2650 atomic->num_components = glsl_get_vector_elements(ptr->type->type);
2651 if (ptr->mode == vtn_variable_mode_ssbo)
2652 atomic->src[src++] = nir_src_for_ssa(index);
2653 atomic->src[src++] = nir_src_for_ssa(offset);
2654 break;
2655
2656 case SpvOpAtomicStore:
2657 atomic->num_components = glsl_get_vector_elements(ptr->type->type);
2658 nir_intrinsic_set_write_mask(atomic, (1 << atomic->num_components) - 1);
2659 atomic->src[src++] = nir_src_for_ssa(vtn_ssa_value(b, w[4])->def);
2660 if (ptr->mode == vtn_variable_mode_ssbo)
2661 atomic->src[src++] = nir_src_for_ssa(index);
2662 atomic->src[src++] = nir_src_for_ssa(offset);
2663 break;
2664
2665 case SpvOpAtomicExchange:
2666 case SpvOpAtomicCompareExchange:
2667 case SpvOpAtomicCompareExchangeWeak:
2668 case SpvOpAtomicIIncrement:
2669 case SpvOpAtomicIDecrement:
2670 case SpvOpAtomicIAdd:
2671 case SpvOpAtomicISub:
2672 case SpvOpAtomicSMin:
2673 case SpvOpAtomicUMin:
2674 case SpvOpAtomicSMax:
2675 case SpvOpAtomicUMax:
2676 case SpvOpAtomicAnd:
2677 case SpvOpAtomicOr:
2678 case SpvOpAtomicXor:
2679 if (ptr->mode == vtn_variable_mode_ssbo)
2680 atomic->src[src++] = nir_src_for_ssa(index);
2681 atomic->src[src++] = nir_src_for_ssa(offset);
2682 fill_common_atomic_sources(b, opcode, w, &atomic->src[src]);
2683 break;
2684
2685 default:
2686 vtn_fail("Invalid SPIR-V atomic");
2687 }
2688 }
2689
2690 if (opcode != SpvOpAtomicStore) {
2691 struct vtn_type *type = vtn_value(b, w[1], vtn_value_type_type)->type;
2692
2693 nir_ssa_dest_init(&atomic->instr, &atomic->dest,
2694 glsl_get_vector_elements(type->type),
2695 glsl_get_bit_size(type->type), NULL);
2696
2697 struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_ssa);
2698 val->ssa = rzalloc(b, struct vtn_ssa_value);
2699 val->ssa->def = &atomic->dest.ssa;
2700 val->ssa->type = type->type;
2701 }
2702
2703 nir_builder_instr_insert(&b->nb, &atomic->instr);
2704 }
2705
2706 static nir_alu_instr *
2707 create_vec(struct vtn_builder *b, unsigned num_components, unsigned bit_size)
2708 {
2709 nir_op op;
2710 switch (num_components) {
2711 case 1: op = nir_op_fmov; break;
2712 case 2: op = nir_op_vec2; break;
2713 case 3: op = nir_op_vec3; break;
2714 case 4: op = nir_op_vec4; break;
2715 default: vtn_fail("bad vector size");
2716 }
2717
2718 nir_alu_instr *vec = nir_alu_instr_create(b->shader, op);
2719 nir_ssa_dest_init(&vec->instr, &vec->dest.dest, num_components,
2720 bit_size, NULL);
2721 vec->dest.write_mask = (1 << num_components) - 1;
2722
2723 return vec;
2724 }
2725
2726 struct vtn_ssa_value *
2727 vtn_ssa_transpose(struct vtn_builder *b, struct vtn_ssa_value *src)
2728 {
2729 if (src->transposed)
2730 return src->transposed;
2731
2732 struct vtn_ssa_value *dest =
2733 vtn_create_ssa_value(b, glsl_transposed_type(src->type));
2734
2735 for (unsigned i = 0; i < glsl_get_matrix_columns(dest->type); i++) {
2736 nir_alu_instr *vec = create_vec(b, glsl_get_matrix_columns(src->type),
2737 glsl_get_bit_size(src->type));
2738 if (glsl_type_is_vector_or_scalar(src->type)) {
2739 vec->src[0].src = nir_src_for_ssa(src->def);
2740 vec->src[0].swizzle[0] = i;
2741 } else {
2742 for (unsigned j = 0; j < glsl_get_matrix_columns(src->type); j++) {
2743 vec->src[j].src = nir_src_for_ssa(src->elems[j]->def);
2744 vec->src[j].swizzle[0] = i;
2745 }
2746 }
2747 nir_builder_instr_insert(&b->nb, &vec->instr);
2748 dest->elems[i]->def = &vec->dest.dest.ssa;
2749 }
2750
2751 dest->transposed = src;
2752
2753 return dest;
2754 }
2755
2756 nir_ssa_def *
2757 vtn_vector_extract(struct vtn_builder *b, nir_ssa_def *src, unsigned index)
2758 {
2759 unsigned swiz[4] = { index };
2760 return nir_swizzle(&b->nb, src, swiz, 1, true);
2761 }
2762
2763 nir_ssa_def *
2764 vtn_vector_insert(struct vtn_builder *b, nir_ssa_def *src, nir_ssa_def *insert,
2765 unsigned index)
2766 {
2767 nir_alu_instr *vec = create_vec(b, src->num_components,
2768 src->bit_size);
2769
2770 for (unsigned i = 0; i < src->num_components; i++) {
2771 if (i == index) {
2772 vec->src[i].src = nir_src_for_ssa(insert);
2773 } else {
2774 vec->src[i].src = nir_src_for_ssa(src);
2775 vec->src[i].swizzle[0] = i;
2776 }
2777 }
2778
2779 nir_builder_instr_insert(&b->nb, &vec->instr);
2780
2781 return &vec->dest.dest.ssa;
2782 }
2783
2784 nir_ssa_def *
2785 vtn_vector_extract_dynamic(struct vtn_builder *b, nir_ssa_def *src,
2786 nir_ssa_def *index)
2787 {
2788 nir_ssa_def *dest = vtn_vector_extract(b, src, 0);
2789 for (unsigned i = 1; i < src->num_components; i++)
2790 dest = nir_bcsel(&b->nb, nir_ieq(&b->nb, index, nir_imm_int(&b->nb, i)),
2791 vtn_vector_extract(b, src, i), dest);
2792
2793 return dest;
2794 }
2795
2796 nir_ssa_def *
2797 vtn_vector_insert_dynamic(struct vtn_builder *b, nir_ssa_def *src,
2798 nir_ssa_def *insert, nir_ssa_def *index)
2799 {
2800 nir_ssa_def *dest = vtn_vector_insert(b, src, insert, 0);
2801 for (unsigned i = 1; i < src->num_components; i++)
2802 dest = nir_bcsel(&b->nb, nir_ieq(&b->nb, index, nir_imm_int(&b->nb, i)),
2803 vtn_vector_insert(b, src, insert, i), dest);
2804
2805 return dest;
2806 }
2807
2808 static nir_ssa_def *
2809 vtn_vector_shuffle(struct vtn_builder *b, unsigned num_components,
2810 nir_ssa_def *src0, nir_ssa_def *src1,
2811 const uint32_t *indices)
2812 {
2813 nir_alu_instr *vec = create_vec(b, num_components, src0->bit_size);
2814
2815 for (unsigned i = 0; i < num_components; i++) {
2816 uint32_t index = indices[i];
2817 if (index == 0xffffffff) {
2818 vec->src[i].src =
2819 nir_src_for_ssa(nir_ssa_undef(&b->nb, 1, src0->bit_size));
2820 } else if (index < src0->num_components) {
2821 vec->src[i].src = nir_src_for_ssa(src0);
2822 vec->src[i].swizzle[0] = index;
2823 } else {
2824 vec->src[i].src = nir_src_for_ssa(src1);
2825 vec->src[i].swizzle[0] = index - src0->num_components;
2826 }
2827 }
2828
2829 nir_builder_instr_insert(&b->nb, &vec->instr);
2830
2831 return &vec->dest.dest.ssa;
2832 }
2833
2834 /*
2835 * Concatentates a number of vectors/scalars together to produce a vector
2836 */
2837 static nir_ssa_def *
2838 vtn_vector_construct(struct vtn_builder *b, unsigned num_components,
2839 unsigned num_srcs, nir_ssa_def **srcs)
2840 {
2841 nir_alu_instr *vec = create_vec(b, num_components, srcs[0]->bit_size);
2842
2843 /* From the SPIR-V 1.1 spec for OpCompositeConstruct:
2844 *
2845 * "When constructing a vector, there must be at least two Constituent
2846 * operands."
2847 */
2848 vtn_assert(num_srcs >= 2);
2849
2850 unsigned dest_idx = 0;
2851 for (unsigned i = 0; i < num_srcs; i++) {
2852 nir_ssa_def *src = srcs[i];
2853 vtn_assert(dest_idx + src->num_components <= num_components);
2854 for (unsigned j = 0; j < src->num_components; j++) {
2855 vec->src[dest_idx].src = nir_src_for_ssa(src);
2856 vec->src[dest_idx].swizzle[0] = j;
2857 dest_idx++;
2858 }
2859 }
2860
2861 /* From the SPIR-V 1.1 spec for OpCompositeConstruct:
2862 *
2863 * "When constructing a vector, the total number of components in all
2864 * the operands must equal the number of components in Result Type."
2865 */
2866 vtn_assert(dest_idx == num_components);
2867
2868 nir_builder_instr_insert(&b->nb, &vec->instr);
2869
2870 return &vec->dest.dest.ssa;
2871 }
2872
2873 static struct vtn_ssa_value *
2874 vtn_composite_copy(void *mem_ctx, struct vtn_ssa_value *src)
2875 {
2876 struct vtn_ssa_value *dest = rzalloc(mem_ctx, struct vtn_ssa_value);
2877 dest->type = src->type;
2878
2879 if (glsl_type_is_vector_or_scalar(src->type)) {
2880 dest->def = src->def;
2881 } else {
2882 unsigned elems = glsl_get_length(src->type);
2883
2884 dest->elems = ralloc_array(mem_ctx, struct vtn_ssa_value *, elems);
2885 for (unsigned i = 0; i < elems; i++)
2886 dest->elems[i] = vtn_composite_copy(mem_ctx, src->elems[i]);
2887 }
2888
2889 return dest;
2890 }
2891
2892 static struct vtn_ssa_value *
2893 vtn_composite_insert(struct vtn_builder *b, struct vtn_ssa_value *src,
2894 struct vtn_ssa_value *insert, const uint32_t *indices,
2895 unsigned num_indices)
2896 {
2897 struct vtn_ssa_value *dest = vtn_composite_copy(b, src);
2898
2899 struct vtn_ssa_value *cur = dest;
2900 unsigned i;
2901 for (i = 0; i < num_indices - 1; i++) {
2902 cur = cur->elems[indices[i]];
2903 }
2904
2905 if (glsl_type_is_vector_or_scalar(cur->type)) {
2906 /* According to the SPIR-V spec, OpCompositeInsert may work down to
2907 * the component granularity. In that case, the last index will be
2908 * the index to insert the scalar into the vector.
2909 */
2910
2911 cur->def = vtn_vector_insert(b, cur->def, insert->def, indices[i]);
2912 } else {
2913 cur->elems[indices[i]] = insert;
2914 }
2915
2916 return dest;
2917 }
2918
2919 static struct vtn_ssa_value *
2920 vtn_composite_extract(struct vtn_builder *b, struct vtn_ssa_value *src,
2921 const uint32_t *indices, unsigned num_indices)
2922 {
2923 struct vtn_ssa_value *cur = src;
2924 for (unsigned i = 0; i < num_indices; i++) {
2925 if (glsl_type_is_vector_or_scalar(cur->type)) {
2926 vtn_assert(i == num_indices - 1);
2927 /* According to the SPIR-V spec, OpCompositeExtract may work down to
2928 * the component granularity. The last index will be the index of the
2929 * vector to extract.
2930 */
2931
2932 struct vtn_ssa_value *ret = rzalloc(b, struct vtn_ssa_value);
2933 ret->type = glsl_scalar_type(glsl_get_base_type(cur->type));
2934 ret->def = vtn_vector_extract(b, cur->def, indices[i]);
2935 return ret;
2936 } else {
2937 cur = cur->elems[indices[i]];
2938 }
2939 }
2940
2941 return cur;
2942 }
2943
2944 static void
2945 vtn_handle_composite(struct vtn_builder *b, SpvOp opcode,
2946 const uint32_t *w, unsigned count)
2947 {
2948 struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_ssa);
2949 const struct glsl_type *type =
2950 vtn_value(b, w[1], vtn_value_type_type)->type->type;
2951 val->ssa = vtn_create_ssa_value(b, type);
2952
2953 switch (opcode) {
2954 case SpvOpVectorExtractDynamic:
2955 val->ssa->def = vtn_vector_extract_dynamic(b, vtn_ssa_value(b, w[3])->def,
2956 vtn_ssa_value(b, w[4])->def);
2957 break;
2958
2959 case SpvOpVectorInsertDynamic:
2960 val->ssa->def = vtn_vector_insert_dynamic(b, vtn_ssa_value(b, w[3])->def,
2961 vtn_ssa_value(b, w[4])->def,
2962 vtn_ssa_value(b, w[5])->def);
2963 break;
2964
2965 case SpvOpVectorShuffle:
2966 val->ssa->def = vtn_vector_shuffle(b, glsl_get_vector_elements(type),
2967 vtn_ssa_value(b, w[3])->def,
2968 vtn_ssa_value(b, w[4])->def,
2969 w + 5);
2970 break;
2971
2972 case SpvOpCompositeConstruct: {
2973 unsigned elems = count - 3;
2974 assume(elems >= 1);
2975 if (glsl_type_is_vector_or_scalar(type)) {
2976 nir_ssa_def *srcs[4];
2977 for (unsigned i = 0; i < elems; i++)
2978 srcs[i] = vtn_ssa_value(b, w[3 + i])->def;
2979 val->ssa->def =
2980 vtn_vector_construct(b, glsl_get_vector_elements(type),
2981 elems, srcs);
2982 } else {
2983 val->ssa->elems = ralloc_array(b, struct vtn_ssa_value *, elems);
2984 for (unsigned i = 0; i < elems; i++)
2985 val->ssa->elems[i] = vtn_ssa_value(b, w[3 + i]);
2986 }
2987 break;
2988 }
2989 case SpvOpCompositeExtract:
2990 val->ssa = vtn_composite_extract(b, vtn_ssa_value(b, w[3]),
2991 w + 4, count - 4);
2992 break;
2993
2994 case SpvOpCompositeInsert:
2995 val->ssa = vtn_composite_insert(b, vtn_ssa_value(b, w[4]),
2996 vtn_ssa_value(b, w[3]),
2997 w + 5, count - 5);
2998 break;
2999
3000 case SpvOpCopyObject:
3001 val->ssa = vtn_composite_copy(b, vtn_ssa_value(b, w[3]));
3002 break;
3003
3004 default:
3005 vtn_fail("unknown composite operation");
3006 }
3007 }
3008
3009 static void
3010 vtn_emit_barrier(struct vtn_builder *b, nir_intrinsic_op op)
3011 {
3012 nir_intrinsic_instr *intrin = nir_intrinsic_instr_create(b->shader, op);
3013 nir_builder_instr_insert(&b->nb, &intrin->instr);
3014 }
3015
3016 static void
3017 vtn_emit_memory_barrier(struct vtn_builder *b, SpvScope scope,
3018 SpvMemorySemanticsMask semantics)
3019 {
3020 static const SpvMemorySemanticsMask all_memory_semantics =
3021 SpvMemorySemanticsUniformMemoryMask |
3022 SpvMemorySemanticsWorkgroupMemoryMask |
3023 SpvMemorySemanticsAtomicCounterMemoryMask |
3024 SpvMemorySemanticsImageMemoryMask;
3025
3026 /* If we're not actually doing a memory barrier, bail */
3027 if (!(semantics & all_memory_semantics))
3028 return;
3029
3030 /* GL and Vulkan don't have these */
3031 vtn_assert(scope != SpvScopeCrossDevice);
3032
3033 if (scope == SpvScopeSubgroup)
3034 return; /* Nothing to do here */
3035
3036 if (scope == SpvScopeWorkgroup) {
3037 vtn_emit_barrier(b, nir_intrinsic_group_memory_barrier);
3038 return;
3039 }
3040
3041 /* There's only two scopes thing left */
3042 vtn_assert(scope == SpvScopeInvocation || scope == SpvScopeDevice);
3043
3044 if ((semantics & all_memory_semantics) == all_memory_semantics) {
3045 vtn_emit_barrier(b, nir_intrinsic_memory_barrier);
3046 return;
3047 }
3048
3049 /* Issue a bunch of more specific barriers */
3050 uint32_t bits = semantics;
3051 while (bits) {
3052 SpvMemorySemanticsMask semantic = 1 << u_bit_scan(&bits);
3053 switch (semantic) {
3054 case SpvMemorySemanticsUniformMemoryMask:
3055 vtn_emit_barrier(b, nir_intrinsic_memory_barrier_buffer);
3056 break;
3057 case SpvMemorySemanticsWorkgroupMemoryMask:
3058 vtn_emit_barrier(b, nir_intrinsic_memory_barrier_shared);
3059 break;
3060 case SpvMemorySemanticsAtomicCounterMemoryMask:
3061 vtn_emit_barrier(b, nir_intrinsic_memory_barrier_atomic_counter);
3062 break;
3063 case SpvMemorySemanticsImageMemoryMask:
3064 vtn_emit_barrier(b, nir_intrinsic_memory_barrier_image);
3065 break;
3066 default:
3067 break;;
3068 }
3069 }
3070 }
3071
3072 static void
3073 vtn_handle_barrier(struct vtn_builder *b, SpvOp opcode,
3074 const uint32_t *w, unsigned count)
3075 {
3076 switch (opcode) {
3077 case SpvOpEmitVertex:
3078 case SpvOpEmitStreamVertex:
3079 case SpvOpEndPrimitive:
3080 case SpvOpEndStreamPrimitive: {
3081 nir_intrinsic_op intrinsic_op;
3082 switch (opcode) {
3083 case SpvOpEmitVertex:
3084 case SpvOpEmitStreamVertex:
3085 intrinsic_op = nir_intrinsic_emit_vertex;
3086 break;
3087 case SpvOpEndPrimitive:
3088 case SpvOpEndStreamPrimitive:
3089 intrinsic_op = nir_intrinsic_end_primitive;
3090 break;
3091 default:
3092 unreachable("Invalid opcode");
3093 }
3094
3095 nir_intrinsic_instr *intrin =
3096 nir_intrinsic_instr_create(b->shader, intrinsic_op);
3097
3098 switch (opcode) {
3099 case SpvOpEmitStreamVertex:
3100 case SpvOpEndStreamPrimitive:
3101 nir_intrinsic_set_stream_id(intrin, w[1]);
3102 break;
3103 default:
3104 break;
3105 }
3106
3107 nir_builder_instr_insert(&b->nb, &intrin->instr);
3108 break;
3109 }
3110
3111 case SpvOpMemoryBarrier: {
3112 SpvScope scope = vtn_constant_value(b, w[1])->values[0].u32[0];
3113 SpvMemorySemanticsMask semantics =
3114 vtn_constant_value(b, w[2])->values[0].u32[0];
3115 vtn_emit_memory_barrier(b, scope, semantics);
3116 return;
3117 }
3118
3119 case SpvOpControlBarrier: {
3120 SpvScope execution_scope =
3121 vtn_constant_value(b, w[1])->values[0].u32[0];
3122 if (execution_scope == SpvScopeWorkgroup)
3123 vtn_emit_barrier(b, nir_intrinsic_barrier);
3124
3125 SpvScope memory_scope =
3126 vtn_constant_value(b, w[2])->values[0].u32[0];
3127 SpvMemorySemanticsMask memory_semantics =
3128 vtn_constant_value(b, w[3])->values[0].u32[0];
3129 vtn_emit_memory_barrier(b, memory_scope, memory_semantics);
3130 break;
3131 }
3132
3133 default:
3134 unreachable("unknown barrier instruction");
3135 }
3136 }
3137
3138 static unsigned
3139 gl_primitive_from_spv_execution_mode(struct vtn_builder *b,
3140 SpvExecutionMode mode)
3141 {
3142 switch (mode) {
3143 case SpvExecutionModeInputPoints:
3144 case SpvExecutionModeOutputPoints:
3145 return 0; /* GL_POINTS */
3146 case SpvExecutionModeInputLines:
3147 return 1; /* GL_LINES */
3148 case SpvExecutionModeInputLinesAdjacency:
3149 return 0x000A; /* GL_LINE_STRIP_ADJACENCY_ARB */
3150 case SpvExecutionModeTriangles:
3151 return 4; /* GL_TRIANGLES */
3152 case SpvExecutionModeInputTrianglesAdjacency:
3153 return 0x000C; /* GL_TRIANGLES_ADJACENCY_ARB */
3154 case SpvExecutionModeQuads:
3155 return 7; /* GL_QUADS */
3156 case SpvExecutionModeIsolines:
3157 return 0x8E7A; /* GL_ISOLINES */
3158 case SpvExecutionModeOutputLineStrip:
3159 return 3; /* GL_LINE_STRIP */
3160 case SpvExecutionModeOutputTriangleStrip:
3161 return 5; /* GL_TRIANGLE_STRIP */
3162 default:
3163 vtn_fail("Invalid primitive type");
3164 }
3165 }
3166
3167 static unsigned
3168 vertices_in_from_spv_execution_mode(struct vtn_builder *b,
3169 SpvExecutionMode mode)
3170 {
3171 switch (mode) {
3172 case SpvExecutionModeInputPoints:
3173 return 1;
3174 case SpvExecutionModeInputLines:
3175 return 2;
3176 case SpvExecutionModeInputLinesAdjacency:
3177 return 4;
3178 case SpvExecutionModeTriangles:
3179 return 3;
3180 case SpvExecutionModeInputTrianglesAdjacency:
3181 return 6;
3182 default:
3183 vtn_fail("Invalid GS input mode");
3184 }
3185 }
3186
3187 static gl_shader_stage
3188 stage_for_execution_model(struct vtn_builder *b, SpvExecutionModel model)
3189 {
3190 switch (model) {
3191 case SpvExecutionModelVertex:
3192 return MESA_SHADER_VERTEX;
3193 case SpvExecutionModelTessellationControl:
3194 return MESA_SHADER_TESS_CTRL;
3195 case SpvExecutionModelTessellationEvaluation:
3196 return MESA_SHADER_TESS_EVAL;
3197 case SpvExecutionModelGeometry:
3198 return MESA_SHADER_GEOMETRY;
3199 case SpvExecutionModelFragment:
3200 return MESA_SHADER_FRAGMENT;
3201 case SpvExecutionModelGLCompute:
3202 return MESA_SHADER_COMPUTE;
3203 default:
3204 vtn_fail("Unsupported execution model");
3205 }
3206 }
3207
3208 #define spv_check_supported(name, cap) do { \
3209 if (!(b->options && b->options->caps.name)) \
3210 vtn_warn("Unsupported SPIR-V capability: %s", \
3211 spirv_capability_to_string(cap)); \
3212 } while(0)
3213
3214
3215 void
3216 vtn_handle_entry_point(struct vtn_builder *b, const uint32_t *w,
3217 unsigned count)
3218 {
3219 struct vtn_value *entry_point = &b->values[w[2]];
3220 /* Let this be a name label regardless */
3221 unsigned name_words;
3222 entry_point->name = vtn_string_literal(b, &w[3], count - 3, &name_words);
3223
3224 if (strcmp(entry_point->name, b->entry_point_name) != 0 ||
3225 stage_for_execution_model(b, w[1]) != b->entry_point_stage)
3226 return;
3227
3228 vtn_assert(b->entry_point == NULL);
3229 b->entry_point = entry_point;
3230 }
3231
3232 static bool
3233 vtn_handle_preamble_instruction(struct vtn_builder *b, SpvOp opcode,
3234 const uint32_t *w, unsigned count)
3235 {
3236 switch (opcode) {
3237 case SpvOpSource: {
3238 const char *lang;
3239 switch (w[1]) {
3240 default:
3241 case SpvSourceLanguageUnknown: lang = "unknown"; break;
3242 case SpvSourceLanguageESSL: lang = "ESSL"; break;
3243 case SpvSourceLanguageGLSL: lang = "GLSL"; break;
3244 case SpvSourceLanguageOpenCL_C: lang = "OpenCL C"; break;
3245 case SpvSourceLanguageOpenCL_CPP: lang = "OpenCL C++"; break;
3246 case SpvSourceLanguageHLSL: lang = "HLSL"; break;
3247 }
3248
3249 uint32_t version = w[2];
3250
3251 const char *file =
3252 (count > 3) ? vtn_value(b, w[3], vtn_value_type_string)->str : "";
3253
3254 vtn_info("Parsing SPIR-V from %s %u source file %s", lang, version, file);
3255 break;
3256 }
3257
3258 case SpvOpSourceExtension:
3259 case SpvOpSourceContinued:
3260 case SpvOpExtension:
3261 case SpvOpModuleProcessed:
3262 /* Unhandled, but these are for debug so that's ok. */
3263 break;
3264
3265 case SpvOpCapability: {
3266 SpvCapability cap = w[1];
3267 switch (cap) {
3268 case SpvCapabilityMatrix:
3269 case SpvCapabilityShader:
3270 case SpvCapabilityGeometry:
3271 case SpvCapabilityGeometryPointSize:
3272 case SpvCapabilityUniformBufferArrayDynamicIndexing:
3273 case SpvCapabilitySampledImageArrayDynamicIndexing:
3274 case SpvCapabilityStorageBufferArrayDynamicIndexing:
3275 case SpvCapabilityStorageImageArrayDynamicIndexing:
3276 case SpvCapabilityImageRect:
3277 case SpvCapabilitySampledRect:
3278 case SpvCapabilitySampled1D:
3279 case SpvCapabilityImage1D:
3280 case SpvCapabilitySampledCubeArray:
3281 case SpvCapabilityImageCubeArray:
3282 case SpvCapabilitySampledBuffer:
3283 case SpvCapabilityImageBuffer:
3284 case SpvCapabilityImageQuery:
3285 case SpvCapabilityDerivativeControl:
3286 case SpvCapabilityInterpolationFunction:
3287 case SpvCapabilityMultiViewport:
3288 case SpvCapabilitySampleRateShading:
3289 case SpvCapabilityClipDistance:
3290 case SpvCapabilityCullDistance:
3291 case SpvCapabilityInputAttachment:
3292 case SpvCapabilityImageGatherExtended:
3293 case SpvCapabilityStorageImageExtendedFormats:
3294 break;
3295
3296 case SpvCapabilityGeometryStreams:
3297 case SpvCapabilityLinkage:
3298 case SpvCapabilityVector16:
3299 case SpvCapabilityFloat16Buffer:
3300 case SpvCapabilityFloat16:
3301 case SpvCapabilityInt64Atomics:
3302 case SpvCapabilityAtomicStorage:
3303 case SpvCapabilityStorageImageMultisample:
3304 case SpvCapabilityInt8:
3305 case SpvCapabilitySparseResidency:
3306 case SpvCapabilityMinLod:
3307 case SpvCapabilityTransformFeedback:
3308 vtn_warn("Unsupported SPIR-V capability: %s",
3309 spirv_capability_to_string(cap));
3310 break;
3311
3312 case SpvCapabilityFloat64:
3313 spv_check_supported(float64, cap);
3314 break;
3315 case SpvCapabilityInt64:
3316 spv_check_supported(int64, cap);
3317 break;
3318 case SpvCapabilityInt16:
3319 spv_check_supported(int16, cap);
3320 break;
3321
3322 case SpvCapabilityAddresses:
3323 case SpvCapabilityKernel:
3324 case SpvCapabilityImageBasic:
3325 case SpvCapabilityImageReadWrite:
3326 case SpvCapabilityImageMipmap:
3327 case SpvCapabilityPipes:
3328 case SpvCapabilityGroups:
3329 case SpvCapabilityDeviceEnqueue:
3330 case SpvCapabilityLiteralSampler:
3331 case SpvCapabilityGenericPointer:
3332 vtn_warn("Unsupported OpenCL-style SPIR-V capability: %s",
3333 spirv_capability_to_string(cap));
3334 break;
3335
3336 case SpvCapabilityImageMSArray:
3337 spv_check_supported(image_ms_array, cap);
3338 break;
3339
3340 case SpvCapabilityTessellation:
3341 case SpvCapabilityTessellationPointSize:
3342 spv_check_supported(tessellation, cap);
3343 break;
3344
3345 case SpvCapabilityDrawParameters:
3346 spv_check_supported(draw_parameters, cap);
3347 break;
3348
3349 case SpvCapabilityStorageImageReadWithoutFormat:
3350 spv_check_supported(image_read_without_format, cap);
3351 break;
3352
3353 case SpvCapabilityStorageImageWriteWithoutFormat:
3354 spv_check_supported(image_write_without_format, cap);
3355 break;
3356
3357 case SpvCapabilityDeviceGroup:
3358 spv_check_supported(device_group, cap);
3359 break;
3360
3361 case SpvCapabilityMultiView:
3362 spv_check_supported(multiview, cap);
3363 break;
3364
3365 case SpvCapabilityGroupNonUniform:
3366 spv_check_supported(subgroup_basic, cap);
3367 break;
3368
3369 case SpvCapabilityGroupNonUniformVote:
3370 spv_check_supported(subgroup_vote, cap);
3371 break;
3372
3373 case SpvCapabilitySubgroupBallotKHR:
3374 case SpvCapabilityGroupNonUniformBallot:
3375 spv_check_supported(subgroup_ballot, cap);
3376 break;
3377
3378 case SpvCapabilityGroupNonUniformShuffle:
3379 case SpvCapabilityGroupNonUniformShuffleRelative:
3380 spv_check_supported(subgroup_shuffle, cap);
3381 break;
3382
3383 case SpvCapabilityGroupNonUniformQuad:
3384 spv_check_supported(subgroup_quad, cap);
3385 break;
3386
3387 case SpvCapabilityGroupNonUniformArithmetic:
3388 case SpvCapabilityGroupNonUniformClustered:
3389 spv_check_supported(subgroup_arithmetic, cap);
3390 break;
3391
3392 case SpvCapabilityVariablePointersStorageBuffer:
3393 case SpvCapabilityVariablePointers:
3394 spv_check_supported(variable_pointers, cap);
3395 break;
3396
3397 case SpvCapabilityStorageUniformBufferBlock16:
3398 case SpvCapabilityStorageUniform16:
3399 case SpvCapabilityStoragePushConstant16:
3400 case SpvCapabilityStorageInputOutput16:
3401 spv_check_supported(storage_16bit, cap);
3402 break;
3403
3404 case SpvCapabilityShaderViewportIndexLayerEXT:
3405 spv_check_supported(shader_viewport_index_layer, cap);
3406 break;
3407
3408 case SpvCapabilityInputAttachmentArrayDynamicIndexingEXT:
3409 case SpvCapabilityUniformTexelBufferArrayDynamicIndexingEXT:
3410 case SpvCapabilityStorageTexelBufferArrayDynamicIndexingEXT:
3411 spv_check_supported(descriptor_array_dynamic_indexing, cap);
3412 break;
3413
3414 case SpvCapabilityRuntimeDescriptorArrayEXT:
3415 spv_check_supported(runtime_descriptor_array, cap);
3416 break;
3417
3418 case SpvCapabilityStencilExportEXT:
3419 spv_check_supported(stencil_export, cap);
3420 break;
3421
3422 default:
3423 vtn_fail("Unhandled capability");
3424 }
3425 break;
3426 }
3427
3428 case SpvOpExtInstImport:
3429 vtn_handle_extension(b, opcode, w, count);
3430 break;
3431
3432 case SpvOpMemoryModel:
3433 vtn_assert(w[1] == SpvAddressingModelLogical);
3434 vtn_assert(w[2] == SpvMemoryModelSimple ||
3435 w[2] == SpvMemoryModelGLSL450);
3436 break;
3437
3438 case SpvOpEntryPoint:
3439 vtn_handle_entry_point(b, w, count);
3440 break;
3441
3442 case SpvOpString:
3443 vtn_push_value(b, w[1], vtn_value_type_string)->str =
3444 vtn_string_literal(b, &w[2], count - 2, NULL);
3445 break;
3446
3447 case SpvOpName:
3448 b->values[w[1]].name = vtn_string_literal(b, &w[2], count - 2, NULL);
3449 break;
3450
3451 case SpvOpMemberName:
3452 /* TODO */
3453 break;
3454
3455 case SpvOpExecutionMode:
3456 case SpvOpDecorationGroup:
3457 case SpvOpDecorate:
3458 case SpvOpMemberDecorate:
3459 case SpvOpGroupDecorate:
3460 case SpvOpGroupMemberDecorate:
3461 vtn_handle_decoration(b, opcode, w, count);
3462 break;
3463
3464 default:
3465 return false; /* End of preamble */
3466 }
3467
3468 return true;
3469 }
3470
3471 static void
3472 vtn_handle_execution_mode(struct vtn_builder *b, struct vtn_value *entry_point,
3473 const struct vtn_decoration *mode, void *data)
3474 {
3475 vtn_assert(b->entry_point == entry_point);
3476
3477 switch(mode->exec_mode) {
3478 case SpvExecutionModeOriginUpperLeft:
3479 case SpvExecutionModeOriginLowerLeft:
3480 b->origin_upper_left =
3481 (mode->exec_mode == SpvExecutionModeOriginUpperLeft);
3482 break;
3483
3484 case SpvExecutionModeEarlyFragmentTests:
3485 vtn_assert(b->shader->info.stage == MESA_SHADER_FRAGMENT);
3486 b->shader->info.fs.early_fragment_tests = true;
3487 break;
3488
3489 case SpvExecutionModeInvocations:
3490 vtn_assert(b->shader->info.stage == MESA_SHADER_GEOMETRY);
3491 b->shader->info.gs.invocations = MAX2(1, mode->literals[0]);
3492 break;
3493
3494 case SpvExecutionModeDepthReplacing:
3495 vtn_assert(b->shader->info.stage == MESA_SHADER_FRAGMENT);
3496 b->shader->info.fs.depth_layout = FRAG_DEPTH_LAYOUT_ANY;
3497 break;
3498 case SpvExecutionModeDepthGreater:
3499 vtn_assert(b->shader->info.stage == MESA_SHADER_FRAGMENT);
3500 b->shader->info.fs.depth_layout = FRAG_DEPTH_LAYOUT_GREATER;
3501 break;
3502 case SpvExecutionModeDepthLess:
3503 vtn_assert(b->shader->info.stage == MESA_SHADER_FRAGMENT);
3504 b->shader->info.fs.depth_layout = FRAG_DEPTH_LAYOUT_LESS;
3505 break;
3506 case SpvExecutionModeDepthUnchanged:
3507 vtn_assert(b->shader->info.stage == MESA_SHADER_FRAGMENT);
3508 b->shader->info.fs.depth_layout = FRAG_DEPTH_LAYOUT_UNCHANGED;
3509 break;
3510
3511 case SpvExecutionModeLocalSize:
3512 vtn_assert(b->shader->info.stage == MESA_SHADER_COMPUTE);
3513 b->shader->info.cs.local_size[0] = mode->literals[0];
3514 b->shader->info.cs.local_size[1] = mode->literals[1];
3515 b->shader->info.cs.local_size[2] = mode->literals[2];
3516 break;
3517 case SpvExecutionModeLocalSizeHint:
3518 break; /* Nothing to do with this */
3519
3520 case SpvExecutionModeOutputVertices:
3521 if (b->shader->info.stage == MESA_SHADER_TESS_CTRL ||
3522 b->shader->info.stage == MESA_SHADER_TESS_EVAL) {
3523 b->shader->info.tess.tcs_vertices_out = mode->literals[0];
3524 } else {
3525 vtn_assert(b->shader->info.stage == MESA_SHADER_GEOMETRY);
3526 b->shader->info.gs.vertices_out = mode->literals[0];
3527 }
3528 break;
3529
3530 case SpvExecutionModeInputPoints:
3531 case SpvExecutionModeInputLines:
3532 case SpvExecutionModeInputLinesAdjacency:
3533 case SpvExecutionModeTriangles:
3534 case SpvExecutionModeInputTrianglesAdjacency:
3535 case SpvExecutionModeQuads:
3536 case SpvExecutionModeIsolines:
3537 if (b->shader->info.stage == MESA_SHADER_TESS_CTRL ||
3538 b->shader->info.stage == MESA_SHADER_TESS_EVAL) {
3539 b->shader->info.tess.primitive_mode =
3540 gl_primitive_from_spv_execution_mode(b, mode->exec_mode);
3541 } else {
3542 vtn_assert(b->shader->info.stage == MESA_SHADER_GEOMETRY);
3543 b->shader->info.gs.vertices_in =
3544 vertices_in_from_spv_execution_mode(b, mode->exec_mode);
3545 }
3546 break;
3547
3548 case SpvExecutionModeOutputPoints:
3549 case SpvExecutionModeOutputLineStrip:
3550 case SpvExecutionModeOutputTriangleStrip:
3551 vtn_assert(b->shader->info.stage == MESA_SHADER_GEOMETRY);
3552 b->shader->info.gs.output_primitive =
3553 gl_primitive_from_spv_execution_mode(b, mode->exec_mode);
3554 break;
3555
3556 case SpvExecutionModeSpacingEqual:
3557 vtn_assert(b->shader->info.stage == MESA_SHADER_TESS_CTRL ||
3558 b->shader->info.stage == MESA_SHADER_TESS_EVAL);
3559 b->shader->info.tess.spacing = TESS_SPACING_EQUAL;
3560 break;
3561 case SpvExecutionModeSpacingFractionalEven:
3562 vtn_assert(b->shader->info.stage == MESA_SHADER_TESS_CTRL ||
3563 b->shader->info.stage == MESA_SHADER_TESS_EVAL);
3564 b->shader->info.tess.spacing = TESS_SPACING_FRACTIONAL_EVEN;
3565 break;
3566 case SpvExecutionModeSpacingFractionalOdd:
3567 vtn_assert(b->shader->info.stage == MESA_SHADER_TESS_CTRL ||
3568 b->shader->info.stage == MESA_SHADER_TESS_EVAL);
3569 b->shader->info.tess.spacing = TESS_SPACING_FRACTIONAL_ODD;
3570 break;
3571 case SpvExecutionModeVertexOrderCw:
3572 vtn_assert(b->shader->info.stage == MESA_SHADER_TESS_CTRL ||
3573 b->shader->info.stage == MESA_SHADER_TESS_EVAL);
3574 b->shader->info.tess.ccw = false;
3575 break;
3576 case SpvExecutionModeVertexOrderCcw:
3577 vtn_assert(b->shader->info.stage == MESA_SHADER_TESS_CTRL ||
3578 b->shader->info.stage == MESA_SHADER_TESS_EVAL);
3579 b->shader->info.tess.ccw = true;
3580 break;
3581 case SpvExecutionModePointMode:
3582 vtn_assert(b->shader->info.stage == MESA_SHADER_TESS_CTRL ||
3583 b->shader->info.stage == MESA_SHADER_TESS_EVAL);
3584 b->shader->info.tess.point_mode = true;
3585 break;
3586
3587 case SpvExecutionModePixelCenterInteger:
3588 b->pixel_center_integer = true;
3589 break;
3590
3591 case SpvExecutionModeXfb:
3592 vtn_fail("Unhandled execution mode");
3593 break;
3594
3595 case SpvExecutionModeVecTypeHint:
3596 case SpvExecutionModeContractionOff:
3597 break; /* OpenCL */
3598
3599 case SpvExecutionModeStencilRefReplacingEXT:
3600 vtn_assert(b->shader->info.stage == MESA_SHADER_FRAGMENT);
3601 break;
3602
3603 default:
3604 vtn_fail("Unhandled execution mode");
3605 }
3606 }
3607
3608 static bool
3609 vtn_handle_variable_or_type_instruction(struct vtn_builder *b, SpvOp opcode,
3610 const uint32_t *w, unsigned count)
3611 {
3612 vtn_set_instruction_result_type(b, opcode, w, count);
3613
3614 switch (opcode) {
3615 case SpvOpSource:
3616 case SpvOpSourceContinued:
3617 case SpvOpSourceExtension:
3618 case SpvOpExtension:
3619 case SpvOpCapability:
3620 case SpvOpExtInstImport:
3621 case SpvOpMemoryModel:
3622 case SpvOpEntryPoint:
3623 case SpvOpExecutionMode:
3624 case SpvOpString:
3625 case SpvOpName:
3626 case SpvOpMemberName:
3627 case SpvOpDecorationGroup:
3628 case SpvOpDecorate:
3629 case SpvOpMemberDecorate:
3630 case SpvOpGroupDecorate:
3631 case SpvOpGroupMemberDecorate:
3632 vtn_fail("Invalid opcode types and variables section");
3633 break;
3634
3635 case SpvOpTypeVoid:
3636 case SpvOpTypeBool:
3637 case SpvOpTypeInt:
3638 case SpvOpTypeFloat:
3639 case SpvOpTypeVector:
3640 case SpvOpTypeMatrix:
3641 case SpvOpTypeImage:
3642 case SpvOpTypeSampler:
3643 case SpvOpTypeSampledImage:
3644 case SpvOpTypeArray:
3645 case SpvOpTypeRuntimeArray:
3646 case SpvOpTypeStruct:
3647 case SpvOpTypeOpaque:
3648 case SpvOpTypePointer:
3649 case SpvOpTypeFunction:
3650 case SpvOpTypeEvent:
3651 case SpvOpTypeDeviceEvent:
3652 case SpvOpTypeReserveId:
3653 case SpvOpTypeQueue:
3654 case SpvOpTypePipe:
3655 vtn_handle_type(b, opcode, w, count);
3656 break;
3657
3658 case SpvOpConstantTrue:
3659 case SpvOpConstantFalse:
3660 case SpvOpConstant:
3661 case SpvOpConstantComposite:
3662 case SpvOpConstantSampler:
3663 case SpvOpConstantNull:
3664 case SpvOpSpecConstantTrue:
3665 case SpvOpSpecConstantFalse:
3666 case SpvOpSpecConstant:
3667 case SpvOpSpecConstantComposite:
3668 case SpvOpSpecConstantOp:
3669 vtn_handle_constant(b, opcode, w, count);
3670 break;
3671
3672 case SpvOpUndef:
3673 case SpvOpVariable:
3674 vtn_handle_variables(b, opcode, w, count);
3675 break;
3676
3677 default:
3678 return false; /* End of preamble */
3679 }
3680
3681 return true;
3682 }
3683
3684 static bool
3685 vtn_handle_body_instruction(struct vtn_builder *b, SpvOp opcode,
3686 const uint32_t *w, unsigned count)
3687 {
3688 switch (opcode) {
3689 case SpvOpLabel:
3690 break;
3691
3692 case SpvOpLoopMerge:
3693 case SpvOpSelectionMerge:
3694 /* This is handled by cfg pre-pass and walk_blocks */
3695 break;
3696
3697 case SpvOpUndef: {
3698 struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_undef);
3699 val->type = vtn_value(b, w[1], vtn_value_type_type)->type;
3700 break;
3701 }
3702
3703 case SpvOpExtInst:
3704 vtn_handle_extension(b, opcode, w, count);
3705 break;
3706
3707 case SpvOpVariable:
3708 case SpvOpLoad:
3709 case SpvOpStore:
3710 case SpvOpCopyMemory:
3711 case SpvOpCopyMemorySized:
3712 case SpvOpAccessChain:
3713 case SpvOpPtrAccessChain:
3714 case SpvOpInBoundsAccessChain:
3715 case SpvOpArrayLength:
3716 vtn_handle_variables(b, opcode, w, count);
3717 break;
3718
3719 case SpvOpFunctionCall:
3720 vtn_handle_function_call(b, opcode, w, count);
3721 break;
3722
3723 case SpvOpSampledImage:
3724 case SpvOpImage:
3725 case SpvOpImageSampleImplicitLod:
3726 case SpvOpImageSampleExplicitLod:
3727 case SpvOpImageSampleDrefImplicitLod:
3728 case SpvOpImageSampleDrefExplicitLod:
3729 case SpvOpImageSampleProjImplicitLod:
3730 case SpvOpImageSampleProjExplicitLod:
3731 case SpvOpImageSampleProjDrefImplicitLod:
3732 case SpvOpImageSampleProjDrefExplicitLod:
3733 case SpvOpImageFetch:
3734 case SpvOpImageGather:
3735 case SpvOpImageDrefGather:
3736 case SpvOpImageQuerySizeLod:
3737 case SpvOpImageQueryLod:
3738 case SpvOpImageQueryLevels:
3739 case SpvOpImageQuerySamples:
3740 vtn_handle_texture(b, opcode, w, count);
3741 break;
3742
3743 case SpvOpImageRead:
3744 case SpvOpImageWrite:
3745 case SpvOpImageTexelPointer:
3746 vtn_handle_image(b, opcode, w, count);
3747 break;
3748
3749 case SpvOpImageQuerySize: {
3750 struct vtn_pointer *image =
3751 vtn_value(b, w[3], vtn_value_type_pointer)->pointer;
3752 if (glsl_type_is_image(image->type->type)) {
3753 vtn_handle_image(b, opcode, w, count);
3754 } else {
3755 vtn_assert(glsl_type_is_sampler(image->type->type));
3756 vtn_handle_texture(b, opcode, w, count);
3757 }
3758 break;
3759 }
3760
3761 case SpvOpAtomicLoad:
3762 case SpvOpAtomicExchange:
3763 case SpvOpAtomicCompareExchange:
3764 case SpvOpAtomicCompareExchangeWeak:
3765 case SpvOpAtomicIIncrement:
3766 case SpvOpAtomicIDecrement:
3767 case SpvOpAtomicIAdd:
3768 case SpvOpAtomicISub:
3769 case SpvOpAtomicSMin:
3770 case SpvOpAtomicUMin:
3771 case SpvOpAtomicSMax:
3772 case SpvOpAtomicUMax:
3773 case SpvOpAtomicAnd:
3774 case SpvOpAtomicOr:
3775 case SpvOpAtomicXor: {
3776 struct vtn_value *pointer = vtn_untyped_value(b, w[3]);
3777 if (pointer->value_type == vtn_value_type_image_pointer) {
3778 vtn_handle_image(b, opcode, w, count);
3779 } else {
3780 vtn_assert(pointer->value_type == vtn_value_type_pointer);
3781 vtn_handle_ssbo_or_shared_atomic(b, opcode, w, count);
3782 }
3783 break;
3784 }
3785
3786 case SpvOpAtomicStore: {
3787 struct vtn_value *pointer = vtn_untyped_value(b, w[1]);
3788 if (pointer->value_type == vtn_value_type_image_pointer) {
3789 vtn_handle_image(b, opcode, w, count);
3790 } else {
3791 vtn_assert(pointer->value_type == vtn_value_type_pointer);
3792 vtn_handle_ssbo_or_shared_atomic(b, opcode, w, count);
3793 }
3794 break;
3795 }
3796
3797 case SpvOpSelect: {
3798 /* Handle OpSelect up-front here because it needs to be able to handle
3799 * pointers and not just regular vectors and scalars.
3800 */
3801 struct vtn_value *res_val = vtn_untyped_value(b, w[2]);
3802 struct vtn_value *sel_val = vtn_untyped_value(b, w[3]);
3803 struct vtn_value *obj1_val = vtn_untyped_value(b, w[4]);
3804 struct vtn_value *obj2_val = vtn_untyped_value(b, w[5]);
3805
3806 const struct glsl_type *sel_type;
3807 switch (res_val->type->base_type) {
3808 case vtn_base_type_scalar:
3809 sel_type = glsl_bool_type();
3810 break;
3811 case vtn_base_type_vector:
3812 sel_type = glsl_vector_type(GLSL_TYPE_BOOL, res_val->type->length);
3813 break;
3814 case vtn_base_type_pointer:
3815 /* We need to have actual storage for pointer types */
3816 vtn_fail_if(res_val->type->type == NULL,
3817 "Invalid pointer result type for OpSelect");
3818 sel_type = glsl_bool_type();
3819 break;
3820 default:
3821 vtn_fail("Result type of OpSelect must be a scalar, vector, or pointer");
3822 }
3823
3824 if (unlikely(sel_val->type->type != sel_type)) {
3825 if (sel_val->type->type == glsl_bool_type()) {
3826 /* This case is illegal but some older versions of GLSLang produce
3827 * it. The GLSLang issue was fixed on March 30, 2017:
3828 *
3829 * https://github.com/KhronosGroup/glslang/issues/809
3830 *
3831 * Unfortunately, there are applications in the wild which are
3832 * shipping with this bug so it isn't nice to fail on them so we
3833 * throw a warning instead. It's not actually a problem for us as
3834 * nir_builder will just splat the condition out which is most
3835 * likely what the client wanted anyway.
3836 */
3837 vtn_warn("Condition type of OpSelect must have the same number "
3838 "of components as Result Type");
3839 } else {
3840 vtn_fail("Condition type of OpSelect must be a scalar or vector "
3841 "of Boolean type. It must have the same number of "
3842 "components as Result Type");
3843 }
3844 }
3845
3846 vtn_fail_if(obj1_val->type != res_val->type ||
3847 obj2_val->type != res_val->type,
3848 "Object types must match the result type in OpSelect");
3849
3850 struct vtn_type *res_type = vtn_value(b, w[1], vtn_value_type_type)->type;
3851 struct vtn_ssa_value *ssa = vtn_create_ssa_value(b, res_type->type);
3852 ssa->def = nir_bcsel(&b->nb, vtn_ssa_value(b, w[3])->def,
3853 vtn_ssa_value(b, w[4])->def,
3854 vtn_ssa_value(b, w[5])->def);
3855 vtn_push_ssa(b, w[2], res_type, ssa);
3856 break;
3857 }
3858
3859 case SpvOpSNegate:
3860 case SpvOpFNegate:
3861 case SpvOpNot:
3862 case SpvOpAny:
3863 case SpvOpAll:
3864 case SpvOpConvertFToU:
3865 case SpvOpConvertFToS:
3866 case SpvOpConvertSToF:
3867 case SpvOpConvertUToF:
3868 case SpvOpUConvert:
3869 case SpvOpSConvert:
3870 case SpvOpFConvert:
3871 case SpvOpQuantizeToF16:
3872 case SpvOpConvertPtrToU:
3873 case SpvOpConvertUToPtr:
3874 case SpvOpPtrCastToGeneric:
3875 case SpvOpGenericCastToPtr:
3876 case SpvOpBitcast:
3877 case SpvOpIsNan:
3878 case SpvOpIsInf:
3879 case SpvOpIsFinite:
3880 case SpvOpIsNormal:
3881 case SpvOpSignBitSet:
3882 case SpvOpLessOrGreater:
3883 case SpvOpOrdered:
3884 case SpvOpUnordered:
3885 case SpvOpIAdd:
3886 case SpvOpFAdd:
3887 case SpvOpISub:
3888 case SpvOpFSub:
3889 case SpvOpIMul:
3890 case SpvOpFMul:
3891 case SpvOpUDiv:
3892 case SpvOpSDiv:
3893 case SpvOpFDiv:
3894 case SpvOpUMod:
3895 case SpvOpSRem:
3896 case SpvOpSMod:
3897 case SpvOpFRem:
3898 case SpvOpFMod:
3899 case SpvOpVectorTimesScalar:
3900 case SpvOpDot:
3901 case SpvOpIAddCarry:
3902 case SpvOpISubBorrow:
3903 case SpvOpUMulExtended:
3904 case SpvOpSMulExtended:
3905 case SpvOpShiftRightLogical:
3906 case SpvOpShiftRightArithmetic:
3907 case SpvOpShiftLeftLogical:
3908 case SpvOpLogicalEqual:
3909 case SpvOpLogicalNotEqual:
3910 case SpvOpLogicalOr:
3911 case SpvOpLogicalAnd:
3912 case SpvOpLogicalNot:
3913 case SpvOpBitwiseOr:
3914 case SpvOpBitwiseXor:
3915 case SpvOpBitwiseAnd:
3916 case SpvOpIEqual:
3917 case SpvOpFOrdEqual:
3918 case SpvOpFUnordEqual:
3919 case SpvOpINotEqual:
3920 case SpvOpFOrdNotEqual:
3921 case SpvOpFUnordNotEqual:
3922 case SpvOpULessThan:
3923 case SpvOpSLessThan:
3924 case SpvOpFOrdLessThan:
3925 case SpvOpFUnordLessThan:
3926 case SpvOpUGreaterThan:
3927 case SpvOpSGreaterThan:
3928 case SpvOpFOrdGreaterThan:
3929 case SpvOpFUnordGreaterThan:
3930 case SpvOpULessThanEqual:
3931 case SpvOpSLessThanEqual:
3932 case SpvOpFOrdLessThanEqual:
3933 case SpvOpFUnordLessThanEqual:
3934 case SpvOpUGreaterThanEqual:
3935 case SpvOpSGreaterThanEqual:
3936 case SpvOpFOrdGreaterThanEqual:
3937 case SpvOpFUnordGreaterThanEqual:
3938 case SpvOpDPdx:
3939 case SpvOpDPdy:
3940 case SpvOpFwidth:
3941 case SpvOpDPdxFine:
3942 case SpvOpDPdyFine:
3943 case SpvOpFwidthFine:
3944 case SpvOpDPdxCoarse:
3945 case SpvOpDPdyCoarse:
3946 case SpvOpFwidthCoarse:
3947 case SpvOpBitFieldInsert:
3948 case SpvOpBitFieldSExtract:
3949 case SpvOpBitFieldUExtract:
3950 case SpvOpBitReverse:
3951 case SpvOpBitCount:
3952 case SpvOpTranspose:
3953 case SpvOpOuterProduct:
3954 case SpvOpMatrixTimesScalar:
3955 case SpvOpVectorTimesMatrix:
3956 case SpvOpMatrixTimesVector:
3957 case SpvOpMatrixTimesMatrix:
3958 vtn_handle_alu(b, opcode, w, count);
3959 break;
3960
3961 case SpvOpVectorExtractDynamic:
3962 case SpvOpVectorInsertDynamic:
3963 case SpvOpVectorShuffle:
3964 case SpvOpCompositeConstruct:
3965 case SpvOpCompositeExtract:
3966 case SpvOpCompositeInsert:
3967 case SpvOpCopyObject:
3968 vtn_handle_composite(b, opcode, w, count);
3969 break;
3970
3971 case SpvOpEmitVertex:
3972 case SpvOpEndPrimitive:
3973 case SpvOpEmitStreamVertex:
3974 case SpvOpEndStreamPrimitive:
3975 case SpvOpControlBarrier:
3976 case SpvOpMemoryBarrier:
3977 vtn_handle_barrier(b, opcode, w, count);
3978 break;
3979
3980 case SpvOpGroupNonUniformElect:
3981 case SpvOpGroupNonUniformAll:
3982 case SpvOpGroupNonUniformAny:
3983 case SpvOpGroupNonUniformAllEqual:
3984 case SpvOpGroupNonUniformBroadcast:
3985 case SpvOpGroupNonUniformBroadcastFirst:
3986 case SpvOpGroupNonUniformBallot:
3987 case SpvOpGroupNonUniformInverseBallot:
3988 case SpvOpGroupNonUniformBallotBitExtract:
3989 case SpvOpGroupNonUniformBallotBitCount:
3990 case SpvOpGroupNonUniformBallotFindLSB:
3991 case SpvOpGroupNonUniformBallotFindMSB:
3992 case SpvOpGroupNonUniformShuffle:
3993 case SpvOpGroupNonUniformShuffleXor:
3994 case SpvOpGroupNonUniformShuffleUp:
3995 case SpvOpGroupNonUniformShuffleDown:
3996 case SpvOpGroupNonUniformIAdd:
3997 case SpvOpGroupNonUniformFAdd:
3998 case SpvOpGroupNonUniformIMul:
3999 case SpvOpGroupNonUniformFMul:
4000 case SpvOpGroupNonUniformSMin:
4001 case SpvOpGroupNonUniformUMin:
4002 case SpvOpGroupNonUniformFMin:
4003 case SpvOpGroupNonUniformSMax:
4004 case SpvOpGroupNonUniformUMax:
4005 case SpvOpGroupNonUniformFMax:
4006 case SpvOpGroupNonUniformBitwiseAnd:
4007 case SpvOpGroupNonUniformBitwiseOr:
4008 case SpvOpGroupNonUniformBitwiseXor:
4009 case SpvOpGroupNonUniformLogicalAnd:
4010 case SpvOpGroupNonUniformLogicalOr:
4011 case SpvOpGroupNonUniformLogicalXor:
4012 case SpvOpGroupNonUniformQuadBroadcast:
4013 case SpvOpGroupNonUniformQuadSwap:
4014 vtn_handle_subgroup(b, opcode, w, count);
4015 break;
4016
4017 default:
4018 vtn_fail("Unhandled opcode");
4019 }
4020
4021 return true;
4022 }
4023
4024 struct vtn_builder*
4025 vtn_create_builder(const uint32_t *words, size_t word_count,
4026 gl_shader_stage stage, const char *entry_point_name,
4027 const struct spirv_to_nir_options *options)
4028 {
4029 /* Initialize the vtn_builder object */
4030 struct vtn_builder *b = rzalloc(NULL, struct vtn_builder);
4031 b->spirv = words;
4032 b->spirv_word_count = word_count;
4033 b->file = NULL;
4034 b->line = -1;
4035 b->col = -1;
4036 exec_list_make_empty(&b->functions);
4037 b->entry_point_stage = stage;
4038 b->entry_point_name = entry_point_name;
4039 b->options = options;
4040
4041 /*
4042 * Handle the SPIR-V header (first 5 dwords).
4043 * Can't use vtx_assert() as the setjmp(3) target isn't initialized yet.
4044 */
4045 if (word_count <= 5)
4046 goto fail;
4047
4048 if (words[0] != SpvMagicNumber) {
4049 vtn_err("words[0] was 0x%x, want 0x%x", words[0], SpvMagicNumber);
4050 goto fail;
4051 }
4052 if (words[1] < 0x10000) {
4053 vtn_err("words[1] was 0x%x, want >= 0x10000", words[1]);
4054 goto fail;
4055 }
4056
4057 /* words[2] == generator magic */
4058 unsigned value_id_bound = words[3];
4059 if (words[4] != 0) {
4060 vtn_err("words[4] was %u, want 0", words[4]);
4061 goto fail;
4062 }
4063
4064 b->value_id_bound = value_id_bound;
4065 b->values = rzalloc_array(b, struct vtn_value, value_id_bound);
4066
4067 return b;
4068 fail:
4069 ralloc_free(b);
4070 return NULL;
4071 }
4072
4073 nir_function *
4074 spirv_to_nir(const uint32_t *words, size_t word_count,
4075 struct nir_spirv_specialization *spec, unsigned num_spec,
4076 gl_shader_stage stage, const char *entry_point_name,
4077 const struct spirv_to_nir_options *options,
4078 const nir_shader_compiler_options *nir_options)
4079
4080 {
4081 const uint32_t *word_end = words + word_count;
4082
4083 struct vtn_builder *b = vtn_create_builder(words, word_count,
4084 stage, entry_point_name,
4085 options);
4086
4087 if (b == NULL)
4088 return NULL;
4089
4090 /* See also _vtn_fail() */
4091 if (setjmp(b->fail_jump)) {
4092 ralloc_free(b);
4093 return NULL;
4094 }
4095
4096 /* Skip the SPIR-V header, handled at vtn_create_builder */
4097 words+= 5;
4098
4099 /* Handle all the preamble instructions */
4100 words = vtn_foreach_instruction(b, words, word_end,
4101 vtn_handle_preamble_instruction);
4102
4103 if (b->entry_point == NULL) {
4104 vtn_fail("Entry point not found");
4105 ralloc_free(b);
4106 return NULL;
4107 }
4108
4109 b->shader = nir_shader_create(b, stage, nir_options, NULL);
4110
4111 /* Set shader info defaults */
4112 b->shader->info.gs.invocations = 1;
4113
4114 /* Parse execution modes */
4115 vtn_foreach_execution_mode(b, b->entry_point,
4116 vtn_handle_execution_mode, NULL);
4117
4118 b->specializations = spec;
4119 b->num_specializations = num_spec;
4120
4121 /* Handle all variable, type, and constant instructions */
4122 words = vtn_foreach_instruction(b, words, word_end,
4123 vtn_handle_variable_or_type_instruction);
4124
4125 /* Set types on all vtn_values */
4126 vtn_foreach_instruction(b, words, word_end, vtn_set_instruction_result_type);
4127
4128 vtn_build_cfg(b, words, word_end);
4129
4130 assert(b->entry_point->value_type == vtn_value_type_function);
4131 b->entry_point->func->referenced = true;
4132
4133 bool progress;
4134 do {
4135 progress = false;
4136 foreach_list_typed(struct vtn_function, func, node, &b->functions) {
4137 if (func->referenced && !func->emitted) {
4138 b->const_table = _mesa_hash_table_create(b, _mesa_hash_pointer,
4139 _mesa_key_pointer_equal);
4140
4141 vtn_function_emit(b, func, vtn_handle_body_instruction);
4142 progress = true;
4143 }
4144 }
4145 } while (progress);
4146
4147 /* We sometimes generate bogus derefs that, while never used, give the
4148 * validator a bit of heartburn. Run dead code to get rid of them.
4149 */
4150 nir_opt_dce(b->shader);
4151
4152 vtn_assert(b->entry_point->value_type == vtn_value_type_function);
4153 nir_function *entry_point = b->entry_point->func->impl->function;
4154 vtn_assert(entry_point);
4155
4156 /* Unparent the shader from the vtn_builder before we delete the builder */
4157 ralloc_steal(NULL, b->shader);
4158
4159 ralloc_free(b);
4160
4161 return entry_point;
4162 }