projects / mesa.git / blob
? search:


6825e0d6a82ab8135bef97e4d48f2c30087305b0
[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 unreachable("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 unreachable("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 unreachable("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 unreachable("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 unreachable("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 SpvOpAtomicISub:
2038 case SpvOpAtomicSMin:
2039 case SpvOpAtomicUMin:
2040 case SpvOpAtomicSMax:
2041 case SpvOpAtomicUMax:
2042 case SpvOpAtomicAnd:
2043 case SpvOpAtomicOr:
2044 case SpvOpAtomicXor:
2045 fill_common_atomic_sources(b, opcode, w, &intrin->src[2]);
2046 break;
2047
2048 default:
2049 unreachable("Invalid image opcode");
2050 }
2051
2052 if (opcode != SpvOpImageWrite) {
2053 struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_ssa);
2054 struct vtn_type *type = vtn_value(b, w[1], vtn_value_type_type)->type;
2055
2056 unsigned dest_components =
2057 nir_intrinsic_infos[intrin->intrinsic].dest_components;
2058 if (intrin->intrinsic == nir_intrinsic_image_size) {
2059 dest_components = intrin->num_components =
2060 glsl_get_vector_elements(type->type);
2061 }
2062
2063 nir_ssa_dest_init(&intrin->instr, &intrin->dest,
2064 dest_components, 32, NULL);
2065
2066 nir_builder_instr_insert(&b->nb, &intrin->instr);
2067
2068 val->ssa = vtn_create_ssa_value(b, type->type);
2069 val->ssa->def = &intrin->dest.ssa;
2070 } else {
2071 nir_builder_instr_insert(&b->nb, &intrin->instr);
2072 }
2073 }
2074
2075 static nir_intrinsic_op
2076 get_ssbo_nir_atomic_op(SpvOp opcode)
2077 {
2078 switch (opcode) {
2079 case SpvOpAtomicLoad: return nir_intrinsic_load_ssbo;
2080 case SpvOpAtomicStore: return nir_intrinsic_store_ssbo;
2081 #define OP(S, N) case SpvOp##S: return nir_intrinsic_ssbo_##N;
2082 OP(AtomicExchange, atomic_exchange)
2083 OP(AtomicCompareExchange, atomic_comp_swap)
2084 OP(AtomicIIncrement, atomic_add)
2085 OP(AtomicIDecrement, atomic_add)
2086 OP(AtomicIAdd, atomic_add)
2087 OP(AtomicISub, atomic_add)
2088 OP(AtomicSMin, atomic_imin)
2089 OP(AtomicUMin, atomic_umin)
2090 OP(AtomicSMax, atomic_imax)
2091 OP(AtomicUMax, atomic_umax)
2092 OP(AtomicAnd, atomic_and)
2093 OP(AtomicOr, atomic_or)
2094 OP(AtomicXor, atomic_xor)
2095 #undef OP
2096 default:
2097 unreachable("Invalid SSBO atomic");
2098 }
2099 }
2100
2101 static nir_intrinsic_op
2102 get_shared_nir_atomic_op(SpvOp opcode)
2103 {
2104 switch (opcode) {
2105 case SpvOpAtomicLoad: return nir_intrinsic_load_var;
2106 case SpvOpAtomicStore: return nir_intrinsic_store_var;
2107 #define OP(S, N) case SpvOp##S: return nir_intrinsic_var_##N;
2108 OP(AtomicExchange, atomic_exchange)
2109 OP(AtomicCompareExchange, atomic_comp_swap)
2110 OP(AtomicIIncrement, atomic_add)
2111 OP(AtomicIDecrement, atomic_add)
2112 OP(AtomicIAdd, atomic_add)
2113 OP(AtomicISub, atomic_add)
2114 OP(AtomicSMin, atomic_imin)
2115 OP(AtomicUMin, atomic_umin)
2116 OP(AtomicSMax, atomic_imax)
2117 OP(AtomicUMax, atomic_umax)
2118 OP(AtomicAnd, atomic_and)
2119 OP(AtomicOr, atomic_or)
2120 OP(AtomicXor, atomic_xor)
2121 #undef OP
2122 default:
2123 unreachable("Invalid shared atomic");
2124 }
2125 }
2126
2127 static void
2128 vtn_handle_ssbo_or_shared_atomic(struct vtn_builder *b, SpvOp opcode,
2129 const uint32_t *w, unsigned count)
2130 {
2131 struct vtn_pointer *ptr;
2132 nir_intrinsic_instr *atomic;
2133
2134 switch (opcode) {
2135 case SpvOpAtomicLoad:
2136 case SpvOpAtomicExchange:
2137 case SpvOpAtomicCompareExchange:
2138 case SpvOpAtomicCompareExchangeWeak:
2139 case SpvOpAtomicIIncrement:
2140 case SpvOpAtomicIDecrement:
2141 case SpvOpAtomicIAdd:
2142 case SpvOpAtomicISub:
2143 case SpvOpAtomicSMin:
2144 case SpvOpAtomicUMin:
2145 case SpvOpAtomicSMax:
2146 case SpvOpAtomicUMax:
2147 case SpvOpAtomicAnd:
2148 case SpvOpAtomicOr:
2149 case SpvOpAtomicXor:
2150 ptr = vtn_value(b, w[3], vtn_value_type_pointer)->pointer;
2151 break;
2152
2153 case SpvOpAtomicStore:
2154 ptr = vtn_value(b, w[1], vtn_value_type_pointer)->pointer;
2155 break;
2156
2157 default:
2158 unreachable("Invalid SPIR-V atomic");
2159 }
2160
2161 /*
2162 SpvScope scope = w[4];
2163 SpvMemorySemanticsMask semantics = w[5];
2164 */
2165
2166 if (ptr->mode == vtn_variable_mode_workgroup) {
2167 nir_deref_var *deref = vtn_pointer_to_deref(b, ptr);
2168 const struct glsl_type *deref_type = nir_deref_tail(&deref->deref)->type;
2169 nir_intrinsic_op op = get_shared_nir_atomic_op(opcode);
2170 atomic = nir_intrinsic_instr_create(b->nb.shader, op);
2171 atomic->variables[0] = nir_deref_var_clone(deref, atomic);
2172
2173 switch (opcode) {
2174 case SpvOpAtomicLoad:
2175 atomic->num_components = glsl_get_vector_elements(deref_type);
2176 break;
2177
2178 case SpvOpAtomicStore:
2179 atomic->num_components = glsl_get_vector_elements(deref_type);
2180 nir_intrinsic_set_write_mask(atomic, (1 << atomic->num_components) - 1);
2181 atomic->src[0] = nir_src_for_ssa(vtn_ssa_value(b, w[4])->def);
2182 break;
2183
2184 case SpvOpAtomicExchange:
2185 case SpvOpAtomicCompareExchange:
2186 case SpvOpAtomicCompareExchangeWeak:
2187 case SpvOpAtomicIIncrement:
2188 case SpvOpAtomicIDecrement:
2189 case SpvOpAtomicIAdd:
2190 case SpvOpAtomicISub:
2191 case SpvOpAtomicSMin:
2192 case SpvOpAtomicUMin:
2193 case SpvOpAtomicSMax:
2194 case SpvOpAtomicUMax:
2195 case SpvOpAtomicAnd:
2196 case SpvOpAtomicOr:
2197 case SpvOpAtomicXor:
2198 fill_common_atomic_sources(b, opcode, w, &atomic->src[0]);
2199 break;
2200
2201 default:
2202 unreachable("Invalid SPIR-V atomic");
2203
2204 }
2205 } else {
2206 assert(ptr->mode == vtn_variable_mode_ssbo);
2207 nir_ssa_def *offset, *index;
2208 offset = vtn_pointer_to_offset(b, ptr, &index, NULL);
2209
2210 nir_intrinsic_op op = get_ssbo_nir_atomic_op(opcode);
2211
2212 atomic = nir_intrinsic_instr_create(b->nb.shader, op);
2213
2214 switch (opcode) {
2215 case SpvOpAtomicLoad:
2216 atomic->num_components = glsl_get_vector_elements(ptr->type->type);
2217 atomic->src[0] = nir_src_for_ssa(index);
2218 atomic->src[1] = nir_src_for_ssa(offset);
2219 break;
2220
2221 case SpvOpAtomicStore:
2222 atomic->num_components = glsl_get_vector_elements(ptr->type->type);
2223 nir_intrinsic_set_write_mask(atomic, (1 << atomic->num_components) - 1);
2224 atomic->src[0] = nir_src_for_ssa(vtn_ssa_value(b, w[4])->def);
2225 atomic->src[1] = nir_src_for_ssa(index);
2226 atomic->src[2] = nir_src_for_ssa(offset);
2227 break;
2228
2229 case SpvOpAtomicExchange:
2230 case SpvOpAtomicCompareExchange:
2231 case SpvOpAtomicCompareExchangeWeak:
2232 case SpvOpAtomicIIncrement:
2233 case SpvOpAtomicIDecrement:
2234 case SpvOpAtomicIAdd:
2235 case SpvOpAtomicISub:
2236 case SpvOpAtomicSMin:
2237 case SpvOpAtomicUMin:
2238 case SpvOpAtomicSMax:
2239 case SpvOpAtomicUMax:
2240 case SpvOpAtomicAnd:
2241 case SpvOpAtomicOr:
2242 case SpvOpAtomicXor:
2243 atomic->src[0] = nir_src_for_ssa(index);
2244 atomic->src[1] = nir_src_for_ssa(offset);
2245 fill_common_atomic_sources(b, opcode, w, &atomic->src[2]);
2246 break;
2247
2248 default:
2249 unreachable("Invalid SPIR-V atomic");
2250 }
2251 }
2252
2253 if (opcode != SpvOpAtomicStore) {
2254 struct vtn_type *type = vtn_value(b, w[1], vtn_value_type_type)->type;
2255
2256 nir_ssa_dest_init(&atomic->instr, &atomic->dest,
2257 glsl_get_vector_elements(type->type),
2258 glsl_get_bit_size(type->type), NULL);
2259
2260 struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_ssa);
2261 val->ssa = rzalloc(b, struct vtn_ssa_value);
2262 val->ssa->def = &atomic->dest.ssa;
2263 val->ssa->type = type->type;
2264 }
2265
2266 nir_builder_instr_insert(&b->nb, &atomic->instr);
2267 }
2268
2269 static nir_alu_instr *
2270 create_vec(nir_shader *shader, unsigned num_components, unsigned bit_size)
2271 {
2272 nir_op op;
2273 switch (num_components) {
2274 case 1: op = nir_op_fmov; break;
2275 case 2: op = nir_op_vec2; break;
2276 case 3: op = nir_op_vec3; break;
2277 case 4: op = nir_op_vec4; break;
2278 default: unreachable("bad vector size");
2279 }
2280
2281 nir_alu_instr *vec = nir_alu_instr_create(shader, op);
2282 nir_ssa_dest_init(&vec->instr, &vec->dest.dest, num_components,
2283 bit_size, NULL);
2284 vec->dest.write_mask = (1 << num_components) - 1;
2285
2286 return vec;
2287 }
2288
2289 struct vtn_ssa_value *
2290 vtn_ssa_transpose(struct vtn_builder *b, struct vtn_ssa_value *src)
2291 {
2292 if (src->transposed)
2293 return src->transposed;
2294
2295 struct vtn_ssa_value *dest =
2296 vtn_create_ssa_value(b, glsl_transposed_type(src->type));
2297
2298 for (unsigned i = 0; i < glsl_get_matrix_columns(dest->type); i++) {
2299 nir_alu_instr *vec = create_vec(b->shader,
2300 glsl_get_matrix_columns(src->type),
2301 glsl_get_bit_size(src->type));
2302 if (glsl_type_is_vector_or_scalar(src->type)) {
2303 vec->src[0].src = nir_src_for_ssa(src->def);
2304 vec->src[0].swizzle[0] = i;
2305 } else {
2306 for (unsigned j = 0; j < glsl_get_matrix_columns(src->type); j++) {
2307 vec->src[j].src = nir_src_for_ssa(src->elems[j]->def);
2308 vec->src[j].swizzle[0] = i;
2309 }
2310 }
2311 nir_builder_instr_insert(&b->nb, &vec->instr);
2312 dest->elems[i]->def = &vec->dest.dest.ssa;
2313 }
2314
2315 dest->transposed = src;
2316
2317 return dest;
2318 }
2319
2320 nir_ssa_def *
2321 vtn_vector_extract(struct vtn_builder *b, nir_ssa_def *src, unsigned index)
2322 {
2323 unsigned swiz[4] = { index };
2324 return nir_swizzle(&b->nb, src, swiz, 1, true);
2325 }
2326
2327 nir_ssa_def *
2328 vtn_vector_insert(struct vtn_builder *b, nir_ssa_def *src, nir_ssa_def *insert,
2329 unsigned index)
2330 {
2331 nir_alu_instr *vec = create_vec(b->shader, src->num_components,
2332 src->bit_size);
2333
2334 for (unsigned i = 0; i < src->num_components; i++) {
2335 if (i == index) {
2336 vec->src[i].src = nir_src_for_ssa(insert);
2337 } else {
2338 vec->src[i].src = nir_src_for_ssa(src);
2339 vec->src[i].swizzle[0] = i;
2340 }
2341 }
2342
2343 nir_builder_instr_insert(&b->nb, &vec->instr);
2344
2345 return &vec->dest.dest.ssa;
2346 }
2347
2348 nir_ssa_def *
2349 vtn_vector_extract_dynamic(struct vtn_builder *b, nir_ssa_def *src,
2350 nir_ssa_def *index)
2351 {
2352 nir_ssa_def *dest = vtn_vector_extract(b, src, 0);
2353 for (unsigned i = 1; i < src->num_components; i++)
2354 dest = nir_bcsel(&b->nb, nir_ieq(&b->nb, index, nir_imm_int(&b->nb, i)),
2355 vtn_vector_extract(b, src, i), dest);
2356
2357 return dest;
2358 }
2359
2360 nir_ssa_def *
2361 vtn_vector_insert_dynamic(struct vtn_builder *b, nir_ssa_def *src,
2362 nir_ssa_def *insert, nir_ssa_def *index)
2363 {
2364 nir_ssa_def *dest = vtn_vector_insert(b, src, insert, 0);
2365 for (unsigned i = 1; i < src->num_components; i++)
2366 dest = nir_bcsel(&b->nb, nir_ieq(&b->nb, index, nir_imm_int(&b->nb, i)),
2367 vtn_vector_insert(b, src, insert, i), dest);
2368
2369 return dest;
2370 }
2371
2372 static nir_ssa_def *
2373 vtn_vector_shuffle(struct vtn_builder *b, unsigned num_components,
2374 nir_ssa_def *src0, nir_ssa_def *src1,
2375 const uint32_t *indices)
2376 {
2377 nir_alu_instr *vec = create_vec(b->shader, num_components, src0->bit_size);
2378
2379 for (unsigned i = 0; i < num_components; i++) {
2380 uint32_t index = indices[i];
2381 if (index == 0xffffffff) {
2382 vec->src[i].src =
2383 nir_src_for_ssa(nir_ssa_undef(&b->nb, 1, src0->bit_size));
2384 } else if (index < src0->num_components) {
2385 vec->src[i].src = nir_src_for_ssa(src0);
2386 vec->src[i].swizzle[0] = index;
2387 } else {
2388 vec->src[i].src = nir_src_for_ssa(src1);
2389 vec->src[i].swizzle[0] = index - src0->num_components;
2390 }
2391 }
2392
2393 nir_builder_instr_insert(&b->nb, &vec->instr);
2394
2395 return &vec->dest.dest.ssa;
2396 }
2397
2398 /*
2399 * Concatentates a number of vectors/scalars together to produce a vector
2400 */
2401 static nir_ssa_def *
2402 vtn_vector_construct(struct vtn_builder *b, unsigned num_components,
2403 unsigned num_srcs, nir_ssa_def **srcs)
2404 {
2405 nir_alu_instr *vec = create_vec(b->shader, num_components,
2406 srcs[0]->bit_size);
2407
2408 /* From the SPIR-V 1.1 spec for OpCompositeConstruct:
2409 *
2410 * "When constructing a vector, there must be at least two Constituent
2411 * operands."
2412 */
2413 assert(num_srcs >= 2);
2414
2415 unsigned dest_idx = 0;
2416 for (unsigned i = 0; i < num_srcs; i++) {
2417 nir_ssa_def *src = srcs[i];
2418 assert(dest_idx + src->num_components <= num_components);
2419 for (unsigned j = 0; j < src->num_components; j++) {
2420 vec->src[dest_idx].src = nir_src_for_ssa(src);
2421 vec->src[dest_idx].swizzle[0] = j;
2422 dest_idx++;
2423 }
2424 }
2425
2426 /* From the SPIR-V 1.1 spec for OpCompositeConstruct:
2427 *
2428 * "When constructing a vector, the total number of components in all
2429 * the operands must equal the number of components in Result Type."
2430 */
2431 assert(dest_idx == num_components);
2432
2433 nir_builder_instr_insert(&b->nb, &vec->instr);
2434
2435 return &vec->dest.dest.ssa;
2436 }
2437
2438 static struct vtn_ssa_value *
2439 vtn_composite_copy(void *mem_ctx, struct vtn_ssa_value *src)
2440 {
2441 struct vtn_ssa_value *dest = rzalloc(mem_ctx, struct vtn_ssa_value);
2442 dest->type = src->type;
2443
2444 if (glsl_type_is_vector_or_scalar(src->type)) {
2445 dest->def = src->def;
2446 } else {
2447 unsigned elems = glsl_get_length(src->type);
2448
2449 dest->elems = ralloc_array(mem_ctx, struct vtn_ssa_value *, elems);
2450 for (unsigned i = 0; i < elems; i++)
2451 dest->elems[i] = vtn_composite_copy(mem_ctx, src->elems[i]);
2452 }
2453
2454 return dest;
2455 }
2456
2457 static struct vtn_ssa_value *
2458 vtn_composite_insert(struct vtn_builder *b, struct vtn_ssa_value *src,
2459 struct vtn_ssa_value *insert, const uint32_t *indices,
2460 unsigned num_indices)
2461 {
2462 struct vtn_ssa_value *dest = vtn_composite_copy(b, src);
2463
2464 struct vtn_ssa_value *cur = dest;
2465 unsigned i;
2466 for (i = 0; i < num_indices - 1; i++) {
2467 cur = cur->elems[indices[i]];
2468 }
2469
2470 if (glsl_type_is_vector_or_scalar(cur->type)) {
2471 /* According to the SPIR-V spec, OpCompositeInsert may work down to
2472 * the component granularity. In that case, the last index will be
2473 * the index to insert the scalar into the vector.
2474 */
2475
2476 cur->def = vtn_vector_insert(b, cur->def, insert->def, indices[i]);
2477 } else {
2478 cur->elems[indices[i]] = insert;
2479 }
2480
2481 return dest;
2482 }
2483
2484 static struct vtn_ssa_value *
2485 vtn_composite_extract(struct vtn_builder *b, struct vtn_ssa_value *src,
2486 const uint32_t *indices, unsigned num_indices)
2487 {
2488 struct vtn_ssa_value *cur = src;
2489 for (unsigned i = 0; i < num_indices; i++) {
2490 if (glsl_type_is_vector_or_scalar(cur->type)) {
2491 assert(i == num_indices - 1);
2492 /* According to the SPIR-V spec, OpCompositeExtract may work down to
2493 * the component granularity. The last index will be the index of the
2494 * vector to extract.
2495 */
2496
2497 struct vtn_ssa_value *ret = rzalloc(b, struct vtn_ssa_value);
2498 ret->type = glsl_scalar_type(glsl_get_base_type(cur->type));
2499 ret->def = vtn_vector_extract(b, cur->def, indices[i]);
2500 return ret;
2501 } else {
2502 cur = cur->elems[indices[i]];
2503 }
2504 }
2505
2506 return cur;
2507 }
2508
2509 static void
2510 vtn_handle_composite(struct vtn_builder *b, SpvOp opcode,
2511 const uint32_t *w, unsigned count)
2512 {
2513 struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_ssa);
2514 const struct glsl_type *type =
2515 vtn_value(b, w[1], vtn_value_type_type)->type->type;
2516 val->ssa = vtn_create_ssa_value(b, type);
2517
2518 switch (opcode) {
2519 case SpvOpVectorExtractDynamic:
2520 val->ssa->def = vtn_vector_extract_dynamic(b, vtn_ssa_value(b, w[3])->def,
2521 vtn_ssa_value(b, w[4])->def);
2522 break;
2523
2524 case SpvOpVectorInsertDynamic:
2525 val->ssa->def = vtn_vector_insert_dynamic(b, vtn_ssa_value(b, w[3])->def,
2526 vtn_ssa_value(b, w[4])->def,
2527 vtn_ssa_value(b, w[5])->def);
2528 break;
2529
2530 case SpvOpVectorShuffle:
2531 val->ssa->def = vtn_vector_shuffle(b, glsl_get_vector_elements(type),
2532 vtn_ssa_value(b, w[3])->def,
2533 vtn_ssa_value(b, w[4])->def,
2534 w + 5);
2535 break;
2536
2537 case SpvOpCompositeConstruct: {
2538 unsigned elems = count - 3;
2539 if (glsl_type_is_vector_or_scalar(type)) {
2540 nir_ssa_def *srcs[4];
2541 for (unsigned i = 0; i < elems; i++)
2542 srcs[i] = vtn_ssa_value(b, w[3 + i])->def;
2543 val->ssa->def =
2544 vtn_vector_construct(b, glsl_get_vector_elements(type),
2545 elems, srcs);
2546 } else {
2547 val->ssa->elems = ralloc_array(b, struct vtn_ssa_value *, elems);
2548 for (unsigned i = 0; i < elems; i++)
2549 val->ssa->elems[i] = vtn_ssa_value(b, w[3 + i]);
2550 }
2551 break;
2552 }
2553 case SpvOpCompositeExtract:
2554 val->ssa = vtn_composite_extract(b, vtn_ssa_value(b, w[3]),
2555 w + 4, count - 4);
2556 break;
2557
2558 case SpvOpCompositeInsert:
2559 val->ssa = vtn_composite_insert(b, vtn_ssa_value(b, w[4]),
2560 vtn_ssa_value(b, w[3]),
2561 w + 5, count - 5);
2562 break;
2563
2564 case SpvOpCopyObject:
2565 val->ssa = vtn_composite_copy(b, vtn_ssa_value(b, w[3]));
2566 break;
2567
2568 default:
2569 unreachable("unknown composite operation");
2570 }
2571 }
2572
2573 static void
2574 vtn_handle_barrier(struct vtn_builder *b, SpvOp opcode,
2575 const uint32_t *w, unsigned count)
2576 {
2577 nir_intrinsic_op intrinsic_op;
2578 switch (opcode) {
2579 case SpvOpEmitVertex:
2580 case SpvOpEmitStreamVertex:
2581 intrinsic_op = nir_intrinsic_emit_vertex;
2582 break;
2583 case SpvOpEndPrimitive:
2584 case SpvOpEndStreamPrimitive:
2585 intrinsic_op = nir_intrinsic_end_primitive;
2586 break;
2587 case SpvOpMemoryBarrier:
2588 intrinsic_op = nir_intrinsic_memory_barrier;
2589 break;
2590 case SpvOpControlBarrier:
2591 intrinsic_op = nir_intrinsic_barrier;
2592 break;
2593 default:
2594 unreachable("unknown barrier instruction");
2595 }
2596
2597 nir_intrinsic_instr *intrin =
2598 nir_intrinsic_instr_create(b->shader, intrinsic_op);
2599
2600 if (opcode == SpvOpEmitStreamVertex || opcode == SpvOpEndStreamPrimitive)
2601 nir_intrinsic_set_stream_id(intrin, w[1]);
2602
2603 nir_builder_instr_insert(&b->nb, &intrin->instr);
2604 }
2605
2606 static unsigned
2607 gl_primitive_from_spv_execution_mode(SpvExecutionMode mode)
2608 {
2609 switch (mode) {
2610 case SpvExecutionModeInputPoints:
2611 case SpvExecutionModeOutputPoints:
2612 return 0; /* GL_POINTS */
2613 case SpvExecutionModeInputLines:
2614 return 1; /* GL_LINES */
2615 case SpvExecutionModeInputLinesAdjacency:
2616 return 0x000A; /* GL_LINE_STRIP_ADJACENCY_ARB */
2617 case SpvExecutionModeTriangles:
2618 return 4; /* GL_TRIANGLES */
2619 case SpvExecutionModeInputTrianglesAdjacency:
2620 return 0x000C; /* GL_TRIANGLES_ADJACENCY_ARB */
2621 case SpvExecutionModeQuads:
2622 return 7; /* GL_QUADS */
2623 case SpvExecutionModeIsolines:
2624 return 0x8E7A; /* GL_ISOLINES */
2625 case SpvExecutionModeOutputLineStrip:
2626 return 3; /* GL_LINE_STRIP */
2627 case SpvExecutionModeOutputTriangleStrip:
2628 return 5; /* GL_TRIANGLE_STRIP */
2629 default:
2630 unreachable("Invalid primitive type");
2631 return 4;
2632 }
2633 }
2634
2635 static unsigned
2636 vertices_in_from_spv_execution_mode(SpvExecutionMode mode)
2637 {
2638 switch (mode) {
2639 case SpvExecutionModeInputPoints:
2640 return 1;
2641 case SpvExecutionModeInputLines:
2642 return 2;
2643 case SpvExecutionModeInputLinesAdjacency:
2644 return 4;
2645 case SpvExecutionModeTriangles:
2646 return 3;
2647 case SpvExecutionModeInputTrianglesAdjacency:
2648 return 6;
2649 default:
2650 unreachable("Invalid GS input mode");
2651 return 0;
2652 }
2653 }
2654
2655 static gl_shader_stage
2656 stage_for_execution_model(SpvExecutionModel model)
2657 {
2658 switch (model) {
2659 case SpvExecutionModelVertex:
2660 return MESA_SHADER_VERTEX;
2661 case SpvExecutionModelTessellationControl:
2662 return MESA_SHADER_TESS_CTRL;
2663 case SpvExecutionModelTessellationEvaluation:
2664 return MESA_SHADER_TESS_EVAL;
2665 case SpvExecutionModelGeometry:
2666 return MESA_SHADER_GEOMETRY;
2667 case SpvExecutionModelFragment:
2668 return MESA_SHADER_FRAGMENT;
2669 case SpvExecutionModelGLCompute:
2670 return MESA_SHADER_COMPUTE;
2671 default:
2672 unreachable("Unsupported execution model");
2673 }
2674 }
2675
2676 #define spv_check_supported(name, cap) do { \
2677 if (!(b->ext && b->ext->name)) \
2678 vtn_warn("Unsupported SPIR-V capability: %s", \
2679 spirv_capability_to_string(cap)); \
2680 } while(0)
2681
2682 static bool
2683 vtn_handle_preamble_instruction(struct vtn_builder *b, SpvOp opcode,
2684 const uint32_t *w, unsigned count)
2685 {
2686 switch (opcode) {
2687 case SpvOpSource:
2688 case SpvOpSourceExtension:
2689 case SpvOpSourceContinued:
2690 case SpvOpExtension:
2691 /* Unhandled, but these are for debug so that's ok. */
2692 break;
2693
2694 case SpvOpCapability: {
2695 SpvCapability cap = w[1];
2696 switch (cap) {
2697 case SpvCapabilityMatrix:
2698 case SpvCapabilityShader:
2699 case SpvCapabilityGeometry:
2700 case SpvCapabilityGeometryPointSize:
2701 case SpvCapabilityUniformBufferArrayDynamicIndexing:
2702 case SpvCapabilitySampledImageArrayDynamicIndexing:
2703 case SpvCapabilityStorageBufferArrayDynamicIndexing:
2704 case SpvCapabilityStorageImageArrayDynamicIndexing:
2705 case SpvCapabilityImageRect:
2706 case SpvCapabilitySampledRect:
2707 case SpvCapabilitySampled1D:
2708 case SpvCapabilityImage1D:
2709 case SpvCapabilitySampledCubeArray:
2710 case SpvCapabilityImageCubeArray:
2711 case SpvCapabilitySampledBuffer:
2712 case SpvCapabilityImageBuffer:
2713 case SpvCapabilityImageQuery:
2714 case SpvCapabilityDerivativeControl:
2715 case SpvCapabilityInterpolationFunction:
2716 case SpvCapabilityMultiViewport:
2717 case SpvCapabilitySampleRateShading:
2718 case SpvCapabilityClipDistance:
2719 case SpvCapabilityCullDistance:
2720 case SpvCapabilityInputAttachment:
2721 case SpvCapabilityImageGatherExtended:
2722 case SpvCapabilityStorageImageExtendedFormats:
2723 break;
2724
2725 case SpvCapabilityGeometryStreams:
2726 case SpvCapabilityLinkage:
2727 case SpvCapabilityVector16:
2728 case SpvCapabilityFloat16Buffer:
2729 case SpvCapabilityFloat16:
2730 case SpvCapabilityInt64Atomics:
2731 case SpvCapabilityAtomicStorage:
2732 case SpvCapabilityInt16:
2733 case SpvCapabilityStorageImageMultisample:
2734 case SpvCapabilityInt8:
2735 case SpvCapabilitySparseResidency:
2736 case SpvCapabilityMinLod:
2737 case SpvCapabilityTransformFeedback:
2738 vtn_warn("Unsupported SPIR-V capability: %s",
2739 spirv_capability_to_string(cap));
2740 break;
2741
2742 case SpvCapabilityFloat64:
2743 spv_check_supported(float64, cap);
2744 break;
2745 case SpvCapabilityInt64:
2746 spv_check_supported(int64, cap);
2747 break;
2748
2749 case SpvCapabilityAddresses:
2750 case SpvCapabilityKernel:
2751 case SpvCapabilityImageBasic:
2752 case SpvCapabilityImageReadWrite:
2753 case SpvCapabilityImageMipmap:
2754 case SpvCapabilityPipes:
2755 case SpvCapabilityGroups:
2756 case SpvCapabilityDeviceEnqueue:
2757 case SpvCapabilityLiteralSampler:
2758 case SpvCapabilityGenericPointer:
2759 vtn_warn("Unsupported OpenCL-style SPIR-V capability: %s",
2760 spirv_capability_to_string(cap));
2761 break;
2762
2763 case SpvCapabilityImageMSArray:
2764 spv_check_supported(image_ms_array, cap);
2765 break;
2766
2767 case SpvCapabilityTessellation:
2768 case SpvCapabilityTessellationPointSize:
2769 spv_check_supported(tessellation, cap);
2770 break;
2771
2772 case SpvCapabilityDrawParameters:
2773 spv_check_supported(draw_parameters, cap);
2774 break;
2775
2776 case SpvCapabilityStorageImageReadWithoutFormat:
2777 spv_check_supported(image_read_without_format, cap);
2778 break;
2779
2780 case SpvCapabilityStorageImageWriteWithoutFormat:
2781 spv_check_supported(image_write_without_format, cap);
2782 break;
2783
2784 case SpvCapabilityMultiView:
2785 spv_check_supported(multiview, cap);
2786 break;
2787
2788 case SpvCapabilityVariablePointersStorageBuffer:
2789 case SpvCapabilityVariablePointers:
2790 spv_check_supported(variable_pointers, cap);
2791 break;
2792
2793 default:
2794 unreachable("Unhandled capability");
2795 }
2796 break;
2797 }
2798
2799 case SpvOpExtInstImport:
2800 vtn_handle_extension(b, opcode, w, count);
2801 break;
2802
2803 case SpvOpMemoryModel:
2804 assert(w[1] == SpvAddressingModelLogical);
2805 assert(w[2] == SpvMemoryModelSimple ||
2806 w[2] == SpvMemoryModelGLSL450);
2807 break;
2808
2809 case SpvOpEntryPoint: {
2810 struct vtn_value *entry_point = &b->values[w[2]];
2811 /* Let this be a name label regardless */
2812 unsigned name_words;
2813 entry_point->name = vtn_string_literal(b, &w[3], count - 3, &name_words);
2814
2815 if (strcmp(entry_point->name, b->entry_point_name) != 0 ||
2816 stage_for_execution_model(w[1]) != b->entry_point_stage)
2817 break;
2818
2819 assert(b->entry_point == NULL);
2820 b->entry_point = entry_point;
2821 break;
2822 }
2823
2824 case SpvOpString:
2825 vtn_push_value(b, w[1], vtn_value_type_string)->str =
2826 vtn_string_literal(b, &w[2], count - 2, NULL);
2827 break;
2828
2829 case SpvOpName:
2830 b->values[w[1]].name = vtn_string_literal(b, &w[2], count - 2, NULL);
2831 break;
2832
2833 case SpvOpMemberName:
2834 /* TODO */
2835 break;
2836
2837 case SpvOpExecutionMode:
2838 case SpvOpDecorationGroup:
2839 case SpvOpDecorate:
2840 case SpvOpMemberDecorate:
2841 case SpvOpGroupDecorate:
2842 case SpvOpGroupMemberDecorate:
2843 vtn_handle_decoration(b, opcode, w, count);
2844 break;
2845
2846 default:
2847 return false; /* End of preamble */
2848 }
2849
2850 return true;
2851 }
2852
2853 static void
2854 vtn_handle_execution_mode(struct vtn_builder *b, struct vtn_value *entry_point,
2855 const struct vtn_decoration *mode, void *data)
2856 {
2857 assert(b->entry_point == entry_point);
2858
2859 switch(mode->exec_mode) {
2860 case SpvExecutionModeOriginUpperLeft:
2861 case SpvExecutionModeOriginLowerLeft:
2862 b->origin_upper_left =
2863 (mode->exec_mode == SpvExecutionModeOriginUpperLeft);
2864 break;
2865
2866 case SpvExecutionModeEarlyFragmentTests:
2867 assert(b->shader->info.stage == MESA_SHADER_FRAGMENT);
2868 b->shader->info.fs.early_fragment_tests = true;
2869 break;
2870
2871 case SpvExecutionModeInvocations:
2872 assert(b->shader->info.stage == MESA_SHADER_GEOMETRY);
2873 b->shader->info.gs.invocations = MAX2(1, mode->literals[0]);
2874 break;
2875
2876 case SpvExecutionModeDepthReplacing:
2877 assert(b->shader->info.stage == MESA_SHADER_FRAGMENT);
2878 b->shader->info.fs.depth_layout = FRAG_DEPTH_LAYOUT_ANY;
2879 break;
2880 case SpvExecutionModeDepthGreater:
2881 assert(b->shader->info.stage == MESA_SHADER_FRAGMENT);
2882 b->shader->info.fs.depth_layout = FRAG_DEPTH_LAYOUT_GREATER;
2883 break;
2884 case SpvExecutionModeDepthLess:
2885 assert(b->shader->info.stage == MESA_SHADER_FRAGMENT);
2886 b->shader->info.fs.depth_layout = FRAG_DEPTH_LAYOUT_LESS;
2887 break;
2888 case SpvExecutionModeDepthUnchanged:
2889 assert(b->shader->info.stage == MESA_SHADER_FRAGMENT);
2890 b->shader->info.fs.depth_layout = FRAG_DEPTH_LAYOUT_UNCHANGED;
2891 break;
2892
2893 case SpvExecutionModeLocalSize:
2894 assert(b->shader->info.stage == MESA_SHADER_COMPUTE);
2895 b->shader->info.cs.local_size[0] = mode->literals[0];
2896 b->shader->info.cs.local_size[1] = mode->literals[1];
2897 b->shader->info.cs.local_size[2] = mode->literals[2];
2898 break;
2899 case SpvExecutionModeLocalSizeHint:
2900 break; /* Nothing to do with this */
2901
2902 case SpvExecutionModeOutputVertices:
2903 if (b->shader->info.stage == MESA_SHADER_TESS_CTRL ||
2904 b->shader->info.stage == MESA_SHADER_TESS_EVAL) {
2905 b->shader->info.tess.tcs_vertices_out = mode->literals[0];
2906 } else {
2907 assert(b->shader->info.stage == MESA_SHADER_GEOMETRY);
2908 b->shader->info.gs.vertices_out = mode->literals[0];
2909 }
2910 break;
2911
2912 case SpvExecutionModeInputPoints:
2913 case SpvExecutionModeInputLines:
2914 case SpvExecutionModeInputLinesAdjacency:
2915 case SpvExecutionModeTriangles:
2916 case SpvExecutionModeInputTrianglesAdjacency:
2917 case SpvExecutionModeQuads:
2918 case SpvExecutionModeIsolines:
2919 if (b->shader->info.stage == MESA_SHADER_TESS_CTRL ||
2920 b->shader->info.stage == MESA_SHADER_TESS_EVAL) {
2921 b->shader->info.tess.primitive_mode =
2922 gl_primitive_from_spv_execution_mode(mode->exec_mode);
2923 } else {
2924 assert(b->shader->info.stage == MESA_SHADER_GEOMETRY);
2925 b->shader->info.gs.vertices_in =
2926 vertices_in_from_spv_execution_mode(mode->exec_mode);
2927 }
2928 break;
2929
2930 case SpvExecutionModeOutputPoints:
2931 case SpvExecutionModeOutputLineStrip:
2932 case SpvExecutionModeOutputTriangleStrip:
2933 assert(b->shader->info.stage == MESA_SHADER_GEOMETRY);
2934 b->shader->info.gs.output_primitive =
2935 gl_primitive_from_spv_execution_mode(mode->exec_mode);
2936 break;
2937
2938 case SpvExecutionModeSpacingEqual:
2939 assert(b->shader->info.stage == MESA_SHADER_TESS_CTRL ||
2940 b->shader->info.stage == MESA_SHADER_TESS_EVAL);
2941 b->shader->info.tess.spacing = TESS_SPACING_EQUAL;
2942 break;
2943 case SpvExecutionModeSpacingFractionalEven:
2944 assert(b->shader->info.stage == MESA_SHADER_TESS_CTRL ||
2945 b->shader->info.stage == MESA_SHADER_TESS_EVAL);
2946 b->shader->info.tess.spacing = TESS_SPACING_FRACTIONAL_EVEN;
2947 break;
2948 case SpvExecutionModeSpacingFractionalOdd:
2949 assert(b->shader->info.stage == MESA_SHADER_TESS_CTRL ||
2950 b->shader->info.stage == MESA_SHADER_TESS_EVAL);
2951 b->shader->info.tess.spacing = TESS_SPACING_FRACTIONAL_ODD;
2952 break;
2953 case SpvExecutionModeVertexOrderCw:
2954 assert(b->shader->info.stage == MESA_SHADER_TESS_CTRL ||
2955 b->shader->info.stage == MESA_SHADER_TESS_EVAL);
2956 b->shader->info.tess.ccw = false;
2957 break;
2958 case SpvExecutionModeVertexOrderCcw:
2959 assert(b->shader->info.stage == MESA_SHADER_TESS_CTRL ||
2960 b->shader->info.stage == MESA_SHADER_TESS_EVAL);
2961 b->shader->info.tess.ccw = true;
2962 break;
2963 case SpvExecutionModePointMode:
2964 assert(b->shader->info.stage == MESA_SHADER_TESS_CTRL ||
2965 b->shader->info.stage == MESA_SHADER_TESS_EVAL);
2966 b->shader->info.tess.point_mode = true;
2967 break;
2968
2969 case SpvExecutionModePixelCenterInteger:
2970 b->pixel_center_integer = true;
2971 break;
2972
2973 case SpvExecutionModeXfb:
2974 unreachable("Unhandled execution mode");
2975 break;
2976
2977 case SpvExecutionModeVecTypeHint:
2978 case SpvExecutionModeContractionOff:
2979 break; /* OpenCL */
2980
2981 default:
2982 unreachable("Unhandled execution mode");
2983 }
2984 }
2985
2986 static bool
2987 vtn_handle_variable_or_type_instruction(struct vtn_builder *b, SpvOp opcode,
2988 const uint32_t *w, unsigned count)
2989 {
2990 switch (opcode) {
2991 case SpvOpSource:
2992 case SpvOpSourceContinued:
2993 case SpvOpSourceExtension:
2994 case SpvOpExtension:
2995 case SpvOpCapability:
2996 case SpvOpExtInstImport:
2997 case SpvOpMemoryModel:
2998 case SpvOpEntryPoint:
2999 case SpvOpExecutionMode:
3000 case SpvOpString:
3001 case SpvOpName:
3002 case SpvOpMemberName:
3003 case SpvOpDecorationGroup:
3004 case SpvOpDecorate:
3005 case SpvOpMemberDecorate:
3006 case SpvOpGroupDecorate:
3007 case SpvOpGroupMemberDecorate:
3008 unreachable("Invalid opcode types and variables section");
3009 break;
3010
3011 case SpvOpTypeVoid:
3012 case SpvOpTypeBool:
3013 case SpvOpTypeInt:
3014 case SpvOpTypeFloat:
3015 case SpvOpTypeVector:
3016 case SpvOpTypeMatrix:
3017 case SpvOpTypeImage:
3018 case SpvOpTypeSampler:
3019 case SpvOpTypeSampledImage:
3020 case SpvOpTypeArray:
3021 case SpvOpTypeRuntimeArray:
3022 case SpvOpTypeStruct:
3023 case SpvOpTypeOpaque:
3024 case SpvOpTypePointer:
3025 case SpvOpTypeFunction:
3026 case SpvOpTypeEvent:
3027 case SpvOpTypeDeviceEvent:
3028 case SpvOpTypeReserveId:
3029 case SpvOpTypeQueue:
3030 case SpvOpTypePipe:
3031 vtn_handle_type(b, opcode, w, count);
3032 break;
3033
3034 case SpvOpConstantTrue:
3035 case SpvOpConstantFalse:
3036 case SpvOpConstant:
3037 case SpvOpConstantComposite:
3038 case SpvOpConstantSampler:
3039 case SpvOpConstantNull:
3040 case SpvOpSpecConstantTrue:
3041 case SpvOpSpecConstantFalse:
3042 case SpvOpSpecConstant:
3043 case SpvOpSpecConstantComposite:
3044 case SpvOpSpecConstantOp:
3045 vtn_handle_constant(b, opcode, w, count);
3046 break;
3047
3048 case SpvOpUndef:
3049 case SpvOpVariable:
3050 vtn_handle_variables(b, opcode, w, count);
3051 break;
3052
3053 default:
3054 return false; /* End of preamble */
3055 }
3056
3057 return true;
3058 }
3059
3060 static bool
3061 vtn_handle_body_instruction(struct vtn_builder *b, SpvOp opcode,
3062 const uint32_t *w, unsigned count)
3063 {
3064 switch (opcode) {
3065 case SpvOpLabel:
3066 break;
3067
3068 case SpvOpLoopMerge:
3069 case SpvOpSelectionMerge:
3070 /* This is handled by cfg pre-pass and walk_blocks */
3071 break;
3072
3073 case SpvOpUndef: {
3074 struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_undef);
3075 val->type = vtn_value(b, w[1], vtn_value_type_type)->type;
3076 break;
3077 }
3078
3079 case SpvOpExtInst:
3080 vtn_handle_extension(b, opcode, w, count);
3081 break;
3082
3083 case SpvOpVariable:
3084 case SpvOpLoad:
3085 case SpvOpStore:
3086 case SpvOpCopyMemory:
3087 case SpvOpCopyMemorySized:
3088 case SpvOpAccessChain:
3089 case SpvOpPtrAccessChain:
3090 case SpvOpInBoundsAccessChain:
3091 case SpvOpArrayLength:
3092 vtn_handle_variables(b, opcode, w, count);
3093 break;
3094
3095 case SpvOpFunctionCall:
3096 vtn_handle_function_call(b, opcode, w, count);
3097 break;
3098
3099 case SpvOpSampledImage:
3100 case SpvOpImage:
3101 case SpvOpImageSampleImplicitLod:
3102 case SpvOpImageSampleExplicitLod:
3103 case SpvOpImageSampleDrefImplicitLod:
3104 case SpvOpImageSampleDrefExplicitLod:
3105 case SpvOpImageSampleProjImplicitLod:
3106 case SpvOpImageSampleProjExplicitLod:
3107 case SpvOpImageSampleProjDrefImplicitLod:
3108 case SpvOpImageSampleProjDrefExplicitLod:
3109 case SpvOpImageFetch:
3110 case SpvOpImageGather:
3111 case SpvOpImageDrefGather:
3112 case SpvOpImageQuerySizeLod:
3113 case SpvOpImageQueryLod:
3114 case SpvOpImageQueryLevels:
3115 case SpvOpImageQuerySamples:
3116 vtn_handle_texture(b, opcode, w, count);
3117 break;
3118
3119 case SpvOpImageRead:
3120 case SpvOpImageWrite:
3121 case SpvOpImageTexelPointer:
3122 vtn_handle_image(b, opcode, w, count);
3123 break;
3124
3125 case SpvOpImageQuerySize: {
3126 struct vtn_pointer *image =
3127 vtn_value(b, w[3], vtn_value_type_pointer)->pointer;
3128 if (image->mode == vtn_variable_mode_image) {
3129 vtn_handle_image(b, opcode, w, count);
3130 } else {
3131 assert(image->mode == vtn_variable_mode_sampler);
3132 vtn_handle_texture(b, opcode, w, count);
3133 }
3134 break;
3135 }
3136
3137 case SpvOpAtomicLoad:
3138 case SpvOpAtomicExchange:
3139 case SpvOpAtomicCompareExchange:
3140 case SpvOpAtomicCompareExchangeWeak:
3141 case SpvOpAtomicIIncrement:
3142 case SpvOpAtomicIDecrement:
3143 case SpvOpAtomicIAdd:
3144 case SpvOpAtomicISub:
3145 case SpvOpAtomicSMin:
3146 case SpvOpAtomicUMin:
3147 case SpvOpAtomicSMax:
3148 case SpvOpAtomicUMax:
3149 case SpvOpAtomicAnd:
3150 case SpvOpAtomicOr:
3151 case SpvOpAtomicXor: {
3152 struct vtn_value *pointer = vtn_untyped_value(b, w[3]);
3153 if (pointer->value_type == vtn_value_type_image_pointer) {
3154 vtn_handle_image(b, opcode, w, count);
3155 } else {
3156 assert(pointer->value_type == vtn_value_type_pointer);
3157 vtn_handle_ssbo_or_shared_atomic(b, opcode, w, count);
3158 }
3159 break;
3160 }
3161
3162 case SpvOpAtomicStore: {
3163 struct vtn_value *pointer = vtn_untyped_value(b, w[1]);
3164 if (pointer->value_type == vtn_value_type_image_pointer) {
3165 vtn_handle_image(b, opcode, w, count);
3166 } else {
3167 assert(pointer->value_type == vtn_value_type_pointer);
3168 vtn_handle_ssbo_or_shared_atomic(b, opcode, w, count);
3169 }
3170 break;
3171 }
3172
3173 case SpvOpSelect: {
3174 /* Handle OpSelect up-front here because it needs to be able to handle
3175 * pointers and not just regular vectors and scalars.
3176 */
3177 struct vtn_type *res_type = vtn_value(b, w[1], vtn_value_type_type)->type;
3178 struct vtn_ssa_value *ssa = vtn_create_ssa_value(b, res_type->type);
3179 ssa->def = nir_bcsel(&b->nb, vtn_ssa_value(b, w[3])->def,
3180 vtn_ssa_value(b, w[4])->def,
3181 vtn_ssa_value(b, w[5])->def);
3182 vtn_push_ssa(b, w[2], res_type, ssa);
3183 break;
3184 }
3185
3186 case SpvOpSNegate:
3187 case SpvOpFNegate:
3188 case SpvOpNot:
3189 case SpvOpAny:
3190 case SpvOpAll:
3191 case SpvOpConvertFToU:
3192 case SpvOpConvertFToS:
3193 case SpvOpConvertSToF:
3194 case SpvOpConvertUToF:
3195 case SpvOpUConvert:
3196 case SpvOpSConvert:
3197 case SpvOpFConvert:
3198 case SpvOpQuantizeToF16:
3199 case SpvOpConvertPtrToU:
3200 case SpvOpConvertUToPtr:
3201 case SpvOpPtrCastToGeneric:
3202 case SpvOpGenericCastToPtr:
3203 case SpvOpBitcast:
3204 case SpvOpIsNan:
3205 case SpvOpIsInf:
3206 case SpvOpIsFinite:
3207 case SpvOpIsNormal:
3208 case SpvOpSignBitSet:
3209 case SpvOpLessOrGreater:
3210 case SpvOpOrdered:
3211 case SpvOpUnordered:
3212 case SpvOpIAdd:
3213 case SpvOpFAdd:
3214 case SpvOpISub:
3215 case SpvOpFSub:
3216 case SpvOpIMul:
3217 case SpvOpFMul:
3218 case SpvOpUDiv:
3219 case SpvOpSDiv:
3220 case SpvOpFDiv:
3221 case SpvOpUMod:
3222 case SpvOpSRem:
3223 case SpvOpSMod:
3224 case SpvOpFRem:
3225 case SpvOpFMod:
3226 case SpvOpVectorTimesScalar:
3227 case SpvOpDot:
3228 case SpvOpIAddCarry:
3229 case SpvOpISubBorrow:
3230 case SpvOpUMulExtended:
3231 case SpvOpSMulExtended:
3232 case SpvOpShiftRightLogical:
3233 case SpvOpShiftRightArithmetic:
3234 case SpvOpShiftLeftLogical:
3235 case SpvOpLogicalEqual:
3236 case SpvOpLogicalNotEqual:
3237 case SpvOpLogicalOr:
3238 case SpvOpLogicalAnd:
3239 case SpvOpLogicalNot:
3240 case SpvOpBitwiseOr:
3241 case SpvOpBitwiseXor:
3242 case SpvOpBitwiseAnd:
3243 case SpvOpIEqual:
3244 case SpvOpFOrdEqual:
3245 case SpvOpFUnordEqual:
3246 case SpvOpINotEqual:
3247 case SpvOpFOrdNotEqual:
3248 case SpvOpFUnordNotEqual:
3249 case SpvOpULessThan:
3250 case SpvOpSLessThan:
3251 case SpvOpFOrdLessThan:
3252 case SpvOpFUnordLessThan:
3253 case SpvOpUGreaterThan:
3254 case SpvOpSGreaterThan:
3255 case SpvOpFOrdGreaterThan:
3256 case SpvOpFUnordGreaterThan:
3257 case SpvOpULessThanEqual:
3258 case SpvOpSLessThanEqual:
3259 case SpvOpFOrdLessThanEqual:
3260 case SpvOpFUnordLessThanEqual:
3261 case SpvOpUGreaterThanEqual:
3262 case SpvOpSGreaterThanEqual:
3263 case SpvOpFOrdGreaterThanEqual:
3264 case SpvOpFUnordGreaterThanEqual:
3265 case SpvOpDPdx:
3266 case SpvOpDPdy:
3267 case SpvOpFwidth:
3268 case SpvOpDPdxFine:
3269 case SpvOpDPdyFine:
3270 case SpvOpFwidthFine:
3271 case SpvOpDPdxCoarse:
3272 case SpvOpDPdyCoarse:
3273 case SpvOpFwidthCoarse:
3274 case SpvOpBitFieldInsert:
3275 case SpvOpBitFieldSExtract:
3276 case SpvOpBitFieldUExtract:
3277 case SpvOpBitReverse:
3278 case SpvOpBitCount:
3279 case SpvOpTranspose:
3280 case SpvOpOuterProduct:
3281 case SpvOpMatrixTimesScalar:
3282 case SpvOpVectorTimesMatrix:
3283 case SpvOpMatrixTimesVector:
3284 case SpvOpMatrixTimesMatrix:
3285 vtn_handle_alu(b, opcode, w, count);
3286 break;
3287
3288 case SpvOpVectorExtractDynamic:
3289 case SpvOpVectorInsertDynamic:
3290 case SpvOpVectorShuffle:
3291 case SpvOpCompositeConstruct:
3292 case SpvOpCompositeExtract:
3293 case SpvOpCompositeInsert:
3294 case SpvOpCopyObject:
3295 vtn_handle_composite(b, opcode, w, count);
3296 break;
3297
3298 case SpvOpEmitVertex:
3299 case SpvOpEndPrimitive:
3300 case SpvOpEmitStreamVertex:
3301 case SpvOpEndStreamPrimitive:
3302 case SpvOpControlBarrier:
3303 case SpvOpMemoryBarrier:
3304 vtn_handle_barrier(b, opcode, w, count);
3305 break;
3306
3307 default:
3308 unreachable("Unhandled opcode");
3309 }
3310
3311 return true;
3312 }
3313
3314 nir_function *
3315 spirv_to_nir(const uint32_t *words, size_t word_count,
3316 struct nir_spirv_specialization *spec, unsigned num_spec,
3317 gl_shader_stage stage, const char *entry_point_name,
3318 const struct nir_spirv_supported_extensions *ext,
3319 const nir_shader_compiler_options *options)
3320 {
3321 const uint32_t *word_end = words + word_count;
3322
3323 /* Handle the SPIR-V header (first 4 dwords) */
3324 assert(word_count > 5);
3325
3326 assert(words[0] == SpvMagicNumber);
3327 assert(words[1] >= 0x10000);
3328 /* words[2] == generator magic */
3329 unsigned value_id_bound = words[3];
3330 assert(words[4] == 0);
3331
3332 words+= 5;
3333
3334 /* Initialize the stn_builder object */
3335 struct vtn_builder *b = rzalloc(NULL, struct vtn_builder);
3336 b->value_id_bound = value_id_bound;
3337 b->values = rzalloc_array(b, struct vtn_value, value_id_bound);
3338 exec_list_make_empty(&b->functions);
3339 b->entry_point_stage = stage;
3340 b->entry_point_name = entry_point_name;
3341 b->ext = ext;
3342
3343 /* Handle all the preamble instructions */
3344 words = vtn_foreach_instruction(b, words, word_end,
3345 vtn_handle_preamble_instruction);
3346
3347 if (b->entry_point == NULL) {
3348 assert(!"Entry point not found");
3349 ralloc_free(b);
3350 return NULL;
3351 }
3352
3353 b->shader = nir_shader_create(NULL, stage, options, NULL);
3354
3355 /* Set shader info defaults */
3356 b->shader->info.gs.invocations = 1;
3357
3358 /* Parse execution modes */
3359 vtn_foreach_execution_mode(b, b->entry_point,
3360 vtn_handle_execution_mode, NULL);
3361
3362 b->specializations = spec;
3363 b->num_specializations = num_spec;
3364
3365 /* Handle all variable, type, and constant instructions */
3366 words = vtn_foreach_instruction(b, words, word_end,
3367 vtn_handle_variable_or_type_instruction);
3368
3369 vtn_build_cfg(b, words, word_end);
3370
3371 foreach_list_typed(struct vtn_function, func, node, &b->functions) {
3372 b->impl = func->impl;
3373 b->const_table = _mesa_hash_table_create(b, _mesa_hash_pointer,
3374 _mesa_key_pointer_equal);
3375
3376 vtn_function_emit(b, func, vtn_handle_body_instruction);
3377 }
3378
3379 assert(b->entry_point->value_type == vtn_value_type_function);
3380 nir_function *entry_point = b->entry_point->func->impl->function;
3381 assert(entry_point);
3382
3383 ralloc_free(b);
3384
3385 return entry_point;
3386 }