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