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