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