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