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