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