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