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