projects / mesa.git / blob
? search:
summary | shortlog | log | commit | commitdiff | tree
history | raw | HEAD
nir/spirv: Implement OpPtrAccessChain for buffers
[mesa.git] / src / compiler / spirv / spirv_to_nir.c
1 /*
2 * Copyright © 2015 Intel Corporation
3 *
4 * Permission is hereby granted, free of charge, to any person obtaining a
5 * copy of this software and associated documentation files (the "Software"),
6 * to deal in the Software without restriction, including without limitation
7 * the rights to use, copy, modify, merge, publish, distribute, sublicense,
8 * and/or sell copies of the Software, and to permit persons to whom the
9 * Software is furnished to do so, subject to the following conditions:
10 *
11 * The above copyright notice and this permission notice (including the next
12 * paragraph) shall be included in all copies or substantial portions of the
13 * Software.
14 *
15 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
16 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
17 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
18 * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
19 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
20 * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
21 * IN THE SOFTWARE.
22 *
23 * Authors:
24 * Jason Ekstrand (jason@jlekstrand.net)
25 *
26 */
27
28 #include "vtn_private.h"
29 #include "nir/nir_vla.h"
30 #include "nir/nir_control_flow.h"
31 #include "nir/nir_constant_expressions.h"
32 #include "spirv_info.h"
33
34 struct spec_constant_value {
35 bool is_double;
36 union {
37 uint32_t data32;
38 uint64_t data64;
39 };
40 };
41
42 void
43 _vtn_warn(const char *file, int line, const char *msg, ...)
44 {
45 char *formatted;
46 va_list args;
47
48 va_start(args, msg);
49 formatted = ralloc_vasprintf(NULL, msg, args);
50 va_end(args);
51
52 fprintf(stderr, "%s:%d WARNING: %s\n", file, line, formatted);
53
54 ralloc_free(formatted);
55 }
56
57 static struct vtn_ssa_value *
58 vtn_undef_ssa_value(struct vtn_builder *b, const struct glsl_type *type)
59 {
60 struct vtn_ssa_value *val = rzalloc(b, struct vtn_ssa_value);
61 val->type = type;
62
63 if (glsl_type_is_vector_or_scalar(type)) {
64 unsigned num_components = glsl_get_vector_elements(val->type);
65 unsigned bit_size = glsl_get_bit_size(val->type);
66 val->def = nir_ssa_undef(&b->nb, num_components, bit_size);
67 } else {
68 unsigned elems = glsl_get_length(val->type);
69 val->elems = ralloc_array(b, struct vtn_ssa_value *, elems);
70 if (glsl_type_is_matrix(type)) {
71 const struct glsl_type *elem_type =
72 glsl_vector_type(glsl_get_base_type(type),
73 glsl_get_vector_elements(type));
74
75 for (unsigned i = 0; i < elems; i++)
76 val->elems[i] = vtn_undef_ssa_value(b, elem_type);
77 } else if (glsl_type_is_array(type)) {
78 const struct glsl_type *elem_type = glsl_get_array_element(type);
79 for (unsigned i = 0; i < elems; i++)
80 val->elems[i] = vtn_undef_ssa_value(b, elem_type);
81 } else {
82 for (unsigned i = 0; i < elems; i++) {
83 const struct glsl_type *elem_type = glsl_get_struct_field(type, i);
84 val->elems[i] = vtn_undef_ssa_value(b, elem_type);
85 }
86 }
87 }
88
89 return val;
90 }
91
92 static struct vtn_ssa_value *
93 vtn_const_ssa_value(struct vtn_builder *b, nir_constant *constant,
94 const struct glsl_type *type)
95 {
96 struct hash_entry *entry = _mesa_hash_table_search(b->const_table, constant);
97
98 if (entry)
99 return entry->data;
100
101 struct vtn_ssa_value *val = rzalloc(b, struct vtn_ssa_value);
102 val->type = type;
103
104 switch (glsl_get_base_type(type)) {
105 case GLSL_TYPE_INT:
106 case GLSL_TYPE_UINT:
107 case GLSL_TYPE_INT64:
108 case GLSL_TYPE_UINT64:
109 case GLSL_TYPE_BOOL:
110 case GLSL_TYPE_FLOAT:
111 case GLSL_TYPE_DOUBLE: {
112 int bit_size = glsl_get_bit_size(type);
113 if (glsl_type_is_vector_or_scalar(type)) {
114 unsigned num_components = glsl_get_vector_elements(val->type);
115 nir_load_const_instr *load =
116 nir_load_const_instr_create(b->shader, num_components, bit_size);
117
118 load->value = constant->values[0];
119
120 nir_instr_insert_before_cf_list(&b->impl->body, &load->instr);
121 val->def = &load->def;
122 } else {
123 assert(glsl_type_is_matrix(type));
124 unsigned rows = glsl_get_vector_elements(val->type);
125 unsigned columns = glsl_get_matrix_columns(val->type);
126 val->elems = ralloc_array(b, struct vtn_ssa_value *, columns);
127
128 for (unsigned i = 0; i < columns; i++) {
129 struct vtn_ssa_value *col_val = rzalloc(b, struct vtn_ssa_value);
130 col_val->type = glsl_get_column_type(val->type);
131 nir_load_const_instr *load =
132 nir_load_const_instr_create(b->shader, rows, bit_size);
133
134 load->value = constant->values[i];
135
136 nir_instr_insert_before_cf_list(&b->impl->body, &load->instr);
137 col_val->def = &load->def;
138
139 val->elems[i] = col_val;
140 }
141 }
142 break;
143 }
144
145 case GLSL_TYPE_ARRAY: {
146 unsigned elems = glsl_get_length(val->type);
147 val->elems = ralloc_array(b, struct vtn_ssa_value *, elems);
148 const struct glsl_type *elem_type = glsl_get_array_element(val->type);
149 for (unsigned i = 0; i < elems; i++)
150 val->elems[i] = vtn_const_ssa_value(b, constant->elements[i],
151 elem_type);
152 break;
153 }
154
155 case GLSL_TYPE_STRUCT: {
156 unsigned elems = glsl_get_length(val->type);
157 val->elems = ralloc_array(b, struct vtn_ssa_value *, elems);
158 for (unsigned i = 0; i < elems; i++) {
159 const struct glsl_type *elem_type =
160 glsl_get_struct_field(val->type, i);
161 val->elems[i] = vtn_const_ssa_value(b, constant->elements[i],
162 elem_type);
163 }
164 break;
165 }
166
167 default:
168 unreachable("bad constant type");
169 }
170
171 return val;
172 }
173
174 struct vtn_ssa_value *
175 vtn_ssa_value(struct vtn_builder *b, uint32_t value_id)
176 {
177 struct vtn_value *val = vtn_untyped_value(b, value_id);
178 switch (val->value_type) {
179 case vtn_value_type_undef:
180 return vtn_undef_ssa_value(b, val->type->type);
181
182 case vtn_value_type_constant:
183 return vtn_const_ssa_value(b, val->constant, val->const_type);
184
185 case vtn_value_type_ssa:
186 return val->ssa;
187
188 default:
189 unreachable("Invalid type for an SSA value");
190 }
191 }
192
193 static char *
194 vtn_string_literal(struct vtn_builder *b, const uint32_t *words,
195 unsigned word_count, unsigned *words_used)
196 {
197 char *dup = ralloc_strndup(b, (char *)words, word_count * sizeof(*words));
198 if (words_used) {
199 /* Ammount of space taken by the string (including the null) */
200 unsigned len = strlen(dup) + 1;
201 *words_used = DIV_ROUND_UP(len, sizeof(*words));
202 }
203 return dup;
204 }
205
206 const uint32_t *
207 vtn_foreach_instruction(struct vtn_builder *b, const uint32_t *start,
208 const uint32_t *end, vtn_instruction_handler handler)
209 {
210 b->file = NULL;
211 b->line = -1;
212 b->col = -1;
213
214 const uint32_t *w = start;
215 while (w < end) {
216 SpvOp opcode = w[0] & SpvOpCodeMask;
217 unsigned count = w[0] >> SpvWordCountShift;
218 assert(count >= 1 && w + count <= end);
219
220 switch (opcode) {
221 case SpvOpNop:
222 break; /* Do nothing */
223
224 case SpvOpLine:
225 b->file = vtn_value(b, w[1], vtn_value_type_string)->str;
226 b->line = w[2];
227 b->col = w[3];
228 break;
229
230 case SpvOpNoLine:
231 b->file = NULL;
232 b->line = -1;
233 b->col = -1;
234 break;
235
236 default:
237 if (!handler(b, opcode, w, count))
238 return w;
239 break;
240 }
241
242 w += count;
243 }
244 assert(w == end);
245 return w;
246 }
247
248 static void
249 vtn_handle_extension(struct vtn_builder *b, SpvOp opcode,
250 const uint32_t *w, unsigned count)
251 {
252 switch (opcode) {
253 case SpvOpExtInstImport: {
254 struct vtn_value *val = vtn_push_value(b, w[1], vtn_value_type_extension);
255 if (strcmp((const char *)&w[2], "GLSL.std.450") == 0) {
256 val->ext_handler = vtn_handle_glsl450_instruction;
257 } else {
258 assert(!"Unsupported extension");
259 }
260 break;
261 }
262
263 case SpvOpExtInst: {
264 struct vtn_value *val = vtn_value(b, w[3], vtn_value_type_extension);
265 bool handled = val->ext_handler(b, w[4], w, count);
266 (void)handled;
267 assert(handled);
268 break;
269 }
270
271 default:
272 unreachable("Unhandled opcode");
273 }
274 }
275
276 static void
277 _foreach_decoration_helper(struct vtn_builder *b,
278 struct vtn_value *base_value,
279 int parent_member,
280 struct vtn_value *value,
281 vtn_decoration_foreach_cb cb, void *data)
282 {
283 for (struct vtn_decoration *dec = value->decoration; dec; dec = dec->next) {
284 int member;
285 if (dec->scope == VTN_DEC_DECORATION) {
286 member = parent_member;
287 } else if (dec->scope >= VTN_DEC_STRUCT_MEMBER0) {
288 assert(parent_member == -1);
289 member = dec->scope - VTN_DEC_STRUCT_MEMBER0;
290 } else {
291 /* Not a decoration */
292 continue;
293 }
294
295 if (dec->group) {
296 assert(dec->group->value_type == vtn_value_type_decoration_group);
297 _foreach_decoration_helper(b, base_value, member, dec->group,
298 cb, data);
299 } else {
300 cb(b, base_value, member, dec, data);
301 }
302 }
303 }
304
305 /** Iterates (recursively if needed) over all of the decorations on a value
306 *
307 * This function iterates over all of the decorations applied to a given
308 * value. If it encounters a decoration group, it recurses into the group
309 * and iterates over all of those decorations as well.
310 */
311 void
312 vtn_foreach_decoration(struct vtn_builder *b, struct vtn_value *value,
313 vtn_decoration_foreach_cb cb, void *data)
314 {
315 _foreach_decoration_helper(b, value, -1, value, cb, data);
316 }
317
318 void
319 vtn_foreach_execution_mode(struct vtn_builder *b, struct vtn_value *value,
320 vtn_execution_mode_foreach_cb cb, void *data)
321 {
322 for (struct vtn_decoration *dec = value->decoration; dec; dec = dec->next) {
323 if (dec->scope != VTN_DEC_EXECUTION_MODE)
324 continue;
325
326 assert(dec->group == NULL);
327 cb(b, value, dec, data);
328 }
329 }
330
331 static void
332 vtn_handle_decoration(struct vtn_builder *b, SpvOp opcode,
333 const uint32_t *w, unsigned count)
334 {
335 const uint32_t *w_end = w + count;
336 const uint32_t target = w[1];
337 w += 2;
338
339 switch (opcode) {
340 case SpvOpDecorationGroup:
341 vtn_push_value(b, target, vtn_value_type_decoration_group);
342 break;
343
344 case SpvOpDecorate:
345 case SpvOpMemberDecorate:
346 case SpvOpExecutionMode: {
347 struct vtn_value *val = &b->values[target];
348
349 struct vtn_decoration *dec = rzalloc(b, struct vtn_decoration);
350 switch (opcode) {
351 case SpvOpDecorate:
352 dec->scope = VTN_DEC_DECORATION;
353 break;
354 case SpvOpMemberDecorate:
355 dec->scope = VTN_DEC_STRUCT_MEMBER0 + *(w++);
356 break;
357 case SpvOpExecutionMode:
358 dec->scope = VTN_DEC_EXECUTION_MODE;
359 break;
360 default:
361 unreachable("Invalid decoration opcode");
362 }
363 dec->decoration = *(w++);
364 dec->literals = w;
365
366 /* Link into the list */
367 dec->next = val->decoration;
368 val->decoration = dec;
369 break;
370 }
371
372 case SpvOpGroupMemberDecorate:
373 case SpvOpGroupDecorate: {
374 struct vtn_value *group =
375 vtn_value(b, target, vtn_value_type_decoration_group);
376
377 for (; w < w_end; w++) {
378 struct vtn_value *val = vtn_untyped_value(b, *w);
379 struct vtn_decoration *dec = rzalloc(b, struct vtn_decoration);
380
381 dec->group = group;
382 if (opcode == SpvOpGroupDecorate) {
383 dec->scope = VTN_DEC_DECORATION;
384 } else {
385 dec->scope = VTN_DEC_STRUCT_MEMBER0 + *(++w);
386 }
387
388 /* Link into the list */
389 dec->next = val->decoration;
390 val->decoration = dec;
391 }
392 break;
393 }
394
395 default:
396 unreachable("Unhandled opcode");
397 }
398 }
399
400 struct member_decoration_ctx {
401 unsigned num_fields;
402 struct glsl_struct_field *fields;
403 struct vtn_type *type;
404 };
405
406 /* does a shallow copy of a vtn_type */
407
408 static struct vtn_type *
409 vtn_type_copy(struct vtn_builder *b, struct vtn_type *src)
410 {
411 struct vtn_type *dest = ralloc(b, struct vtn_type);
412 *dest = *src;
413
414 switch (src->base_type) {
415 case vtn_base_type_void:
416 case vtn_base_type_scalar:
417 case vtn_base_type_vector:
418 case vtn_base_type_matrix:
419 case vtn_base_type_array:
420 case vtn_base_type_pointer:
421 case vtn_base_type_image:
422 case vtn_base_type_sampler:
423 /* Nothing more to do */
424 break;
425
426 case vtn_base_type_struct:
427 dest->members = ralloc_array(b, struct vtn_type *, src->length);
428 memcpy(dest->members, src->members,
429 src->length * sizeof(src->members[0]));
430
431 dest->offsets = ralloc_array(b, unsigned, src->length);
432 memcpy(dest->offsets, src->offsets,
433 src->length * sizeof(src->offsets[0]));
434 break;
435
436 case vtn_base_type_function:
437 dest->params = ralloc_array(b, struct vtn_type *, src->length);
438 memcpy(dest->params, src->params, src->length * sizeof(src->params[0]));
439 break;
440 }
441
442 return dest;
443 }
444
445 static struct vtn_type *
446 mutable_matrix_member(struct vtn_builder *b, struct vtn_type *type, int member)
447 {
448 type->members[member] = vtn_type_copy(b, type->members[member]);
449 type = type->members[member];
450
451 /* We may have an array of matrices.... Oh, joy! */
452 while (glsl_type_is_array(type->type)) {
453 type->array_element = vtn_type_copy(b, type->array_element);
454 type = type->array_element;
455 }
456
457 assert(glsl_type_is_matrix(type->type));
458
459 return type;
460 }
461
462 static void
463 struct_member_decoration_cb(struct vtn_builder *b,
464 struct vtn_value *val, int member,
465 const struct vtn_decoration *dec, void *void_ctx)
466 {
467 struct member_decoration_ctx *ctx = void_ctx;
468
469 if (member < 0)
470 return;
471
472 assert(member < ctx->num_fields);
473
474 switch (dec->decoration) {
475 case SpvDecorationNonWritable:
476 case SpvDecorationNonReadable:
477 case SpvDecorationRelaxedPrecision:
478 case SpvDecorationVolatile:
479 case SpvDecorationCoherent:
480 case SpvDecorationUniform:
481 break; /* FIXME: Do nothing with this for now. */
482 case SpvDecorationNoPerspective:
483 ctx->fields[member].interpolation = INTERP_MODE_NOPERSPECTIVE;
484 break;
485 case SpvDecorationFlat:
486 ctx->fields[member].interpolation = INTERP_MODE_FLAT;
487 break;
488 case SpvDecorationCentroid:
489 ctx->fields[member].centroid = true;
490 break;
491 case SpvDecorationSample:
492 ctx->fields[member].sample = true;
493 break;
494 case SpvDecorationStream:
495 /* Vulkan only allows one GS stream */
496 assert(dec->literals[0] == 0);
497 break;
498 case SpvDecorationLocation:
499 ctx->fields[member].location = dec->literals[0];
500 break;
501 case SpvDecorationComponent:
502 break; /* FIXME: What should we do with these? */
503 case SpvDecorationBuiltIn:
504 ctx->type->members[member] = vtn_type_copy(b, ctx->type->members[member]);
505 ctx->type->members[member]->is_builtin = true;
506 ctx->type->members[member]->builtin = dec->literals[0];
507 ctx->type->builtin_block = true;
508 break;
509 case SpvDecorationOffset:
510 ctx->type->offsets[member] = dec->literals[0];
511 break;
512 case SpvDecorationMatrixStride:
513 /* Handled as a second pass */
514 break;
515 case SpvDecorationColMajor:
516 break; /* Nothing to do here. Column-major is the default. */
517 case SpvDecorationRowMajor:
518 mutable_matrix_member(b, ctx->type, member)->row_major = true;
519 break;
520
521 case SpvDecorationPatch:
522 break;
523
524 case SpvDecorationSpecId:
525 case SpvDecorationBlock:
526 case SpvDecorationBufferBlock:
527 case SpvDecorationArrayStride:
528 case SpvDecorationGLSLShared:
529 case SpvDecorationGLSLPacked:
530 case SpvDecorationInvariant:
531 case SpvDecorationRestrict:
532 case SpvDecorationAliased:
533 case SpvDecorationConstant:
534 case SpvDecorationIndex:
535 case SpvDecorationBinding:
536 case SpvDecorationDescriptorSet:
537 case SpvDecorationLinkageAttributes:
538 case SpvDecorationNoContraction:
539 case SpvDecorationInputAttachmentIndex:
540 vtn_warn("Decoration not allowed on struct members: %s",
541 spirv_decoration_to_string(dec->decoration));
542 break;
543
544 case SpvDecorationXfbBuffer:
545 case SpvDecorationXfbStride:
546 vtn_warn("Vulkan does not have transform feedback");
547 break;
548
549 case SpvDecorationCPacked:
550 case SpvDecorationSaturatedConversion:
551 case SpvDecorationFuncParamAttr:
552 case SpvDecorationFPRoundingMode:
553 case SpvDecorationFPFastMathMode:
554 case SpvDecorationAlignment:
555 vtn_warn("Decoration only allowed for CL-style kernels: %s",
556 spirv_decoration_to_string(dec->decoration));
557 break;
558
559 default:
560 unreachable("Unhandled decoration");
561 }
562 }
563
564 /* Matrix strides are handled as a separate pass because we need to know
565 * whether the matrix is row-major or not first.
566 */
567 static void
568 struct_member_matrix_stride_cb(struct vtn_builder *b,
569 struct vtn_value *val, int member,
570 const struct vtn_decoration *dec,
571 void *void_ctx)
572 {
573 if (dec->decoration != SpvDecorationMatrixStride)
574 return;
575 assert(member >= 0);
576
577 struct member_decoration_ctx *ctx = void_ctx;
578
579 struct vtn_type *mat_type = mutable_matrix_member(b, ctx->type, member);
580 if (mat_type->row_major) {
581 mat_type->array_element = vtn_type_copy(b, mat_type->array_element);
582 mat_type->stride = mat_type->array_element->stride;
583 mat_type->array_element->stride = dec->literals[0];
584 } else {
585 assert(mat_type->array_element->stride > 0);
586 mat_type->stride = dec->literals[0];
587 }
588 }
589
590 static void
591 type_decoration_cb(struct vtn_builder *b,
592 struct vtn_value *val, int member,
593 const struct vtn_decoration *dec, void *ctx)
594 {
595 struct vtn_type *type = val->type;
596
597 if (member != -1)
598 return;
599
600 switch (dec->decoration) {
601 case SpvDecorationArrayStride:
602 assert(type->base_type == vtn_base_type_matrix ||
603 type->base_type == vtn_base_type_array ||
604 type->base_type == vtn_base_type_pointer);
605 type->stride = dec->literals[0];
606 break;
607 case SpvDecorationBlock:
608 assert(type->base_type == vtn_base_type_struct);
609 type->block = true;
610 break;
611 case SpvDecorationBufferBlock:
612 assert(type->base_type == vtn_base_type_struct);
613 type->buffer_block = true;
614 break;
615 case SpvDecorationGLSLShared:
616 case SpvDecorationGLSLPacked:
617 /* Ignore these, since we get explicit offsets anyways */
618 break;
619
620 case SpvDecorationRowMajor:
621 case SpvDecorationColMajor:
622 case SpvDecorationMatrixStride:
623 case SpvDecorationBuiltIn:
624 case SpvDecorationNoPerspective:
625 case SpvDecorationFlat:
626 case SpvDecorationPatch:
627 case SpvDecorationCentroid:
628 case SpvDecorationSample:
629 case SpvDecorationVolatile:
630 case SpvDecorationCoherent:
631 case SpvDecorationNonWritable:
632 case SpvDecorationNonReadable:
633 case SpvDecorationUniform:
634 case SpvDecorationStream:
635 case SpvDecorationLocation:
636 case SpvDecorationComponent:
637 case SpvDecorationOffset:
638 case SpvDecorationXfbBuffer:
639 case SpvDecorationXfbStride:
640 vtn_warn("Decoration only allowed for struct members: %s",
641 spirv_decoration_to_string(dec->decoration));
642 break;
643
644 case SpvDecorationRelaxedPrecision:
645 case SpvDecorationSpecId:
646 case SpvDecorationInvariant:
647 case SpvDecorationRestrict:
648 case SpvDecorationAliased:
649 case SpvDecorationConstant:
650 case SpvDecorationIndex:
651 case SpvDecorationBinding:
652 case SpvDecorationDescriptorSet:
653 case SpvDecorationLinkageAttributes:
654 case SpvDecorationNoContraction:
655 case SpvDecorationInputAttachmentIndex:
656 vtn_warn("Decoration not allowed on types: %s",
657 spirv_decoration_to_string(dec->decoration));
658 break;
659
660 case SpvDecorationCPacked:
661 case SpvDecorationSaturatedConversion:
662 case SpvDecorationFuncParamAttr:
663 case SpvDecorationFPRoundingMode:
664 case SpvDecorationFPFastMathMode:
665 case SpvDecorationAlignment:
666 vtn_warn("Decoration only allowed for CL-style kernels: %s",
667 spirv_decoration_to_string(dec->decoration));
668 break;
669
670 default:
671 unreachable("Unhandled decoration");
672 }
673 }
674
675 static unsigned
676 translate_image_format(SpvImageFormat format)
677 {
678 switch (format) {
679 case SpvImageFormatUnknown: return 0; /* GL_NONE */
680 case SpvImageFormatRgba32f: return 0x8814; /* GL_RGBA32F */
681 case SpvImageFormatRgba16f: return 0x881A; /* GL_RGBA16F */
682 case SpvImageFormatR32f: return 0x822E; /* GL_R32F */
683 case SpvImageFormatRgba8: return 0x8058; /* GL_RGBA8 */
684 case SpvImageFormatRgba8Snorm: return 0x8F97; /* GL_RGBA8_SNORM */
685 case SpvImageFormatRg32f: return 0x8230; /* GL_RG32F */
686 case SpvImageFormatRg16f: return 0x822F; /* GL_RG16F */
687 case SpvImageFormatR11fG11fB10f: return 0x8C3A; /* GL_R11F_G11F_B10F */
688 case SpvImageFormatR16f: return 0x822D; /* GL_R16F */
689 case SpvImageFormatRgba16: return 0x805B; /* GL_RGBA16 */
690 case SpvImageFormatRgb10A2: return 0x8059; /* GL_RGB10_A2 */
691 case SpvImageFormatRg16: return 0x822C; /* GL_RG16 */
692 case SpvImageFormatRg8: return 0x822B; /* GL_RG8 */
693 case SpvImageFormatR16: return 0x822A; /* GL_R16 */
694 case SpvImageFormatR8: return 0x8229; /* GL_R8 */
695 case SpvImageFormatRgba16Snorm: return 0x8F9B; /* GL_RGBA16_SNORM */
696 case SpvImageFormatRg16Snorm: return 0x8F99; /* GL_RG16_SNORM */
697 case SpvImageFormatRg8Snorm: return 0x8F95; /* GL_RG8_SNORM */
698 case SpvImageFormatR16Snorm: return 0x8F98; /* GL_R16_SNORM */
699 case SpvImageFormatR8Snorm: return 0x8F94; /* GL_R8_SNORM */
700 case SpvImageFormatRgba32i: return 0x8D82; /* GL_RGBA32I */
701 case SpvImageFormatRgba16i: return 0x8D88; /* GL_RGBA16I */
702 case SpvImageFormatRgba8i: return 0x8D8E; /* GL_RGBA8I */
703 case SpvImageFormatR32i: return 0x8235; /* GL_R32I */
704 case SpvImageFormatRg32i: return 0x823B; /* GL_RG32I */
705 case SpvImageFormatRg16i: return 0x8239; /* GL_RG16I */
706 case SpvImageFormatRg8i: return 0x8237; /* GL_RG8I */
707 case SpvImageFormatR16i: return 0x8233; /* GL_R16I */
708 case SpvImageFormatR8i: return 0x8231; /* GL_R8I */
709 case SpvImageFormatRgba32ui: return 0x8D70; /* GL_RGBA32UI */
710 case SpvImageFormatRgba16ui: return 0x8D76; /* GL_RGBA16UI */
711 case SpvImageFormatRgba8ui: return 0x8D7C; /* GL_RGBA8UI */
712 case SpvImageFormatR32ui: return 0x8236; /* GL_R32UI */
713 case SpvImageFormatRgb10a2ui: return 0x906F; /* GL_RGB10_A2UI */
714 case SpvImageFormatRg32ui: return 0x823C; /* GL_RG32UI */
715 case SpvImageFormatRg16ui: return 0x823A; /* GL_RG16UI */
716 case SpvImageFormatRg8ui: return 0x8238; /* GL_RG8UI */
717 case SpvImageFormatR16ui: return 0x823A; /* GL_RG16UI */
718 case SpvImageFormatR8ui: return 0x8232; /* GL_R8UI */
719 default:
720 assert(!"Invalid image format");
721 return 0;
722 }
723 }
724
725 static void
726 vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
727 const uint32_t *w, unsigned count)
728 {
729 struct vtn_value *val = vtn_push_value(b, w[1], vtn_value_type_type);
730
731 val->type = rzalloc(b, struct vtn_type);
732 val->type->val = val;
733
734 switch (opcode) {
735 case SpvOpTypeVoid:
736 val->type->base_type = vtn_base_type_void;
737 val->type->type = glsl_void_type();
738 break;
739 case SpvOpTypeBool:
740 val->type->base_type = vtn_base_type_scalar;
741 val->type->type = glsl_bool_type();
742 break;
743 case SpvOpTypeInt: {
744 int bit_size = w[2];
745 const bool signedness = w[3];
746 val->type->base_type = vtn_base_type_scalar;
747 if (bit_size == 64)
748 val->type->type = (signedness ? glsl_int64_t_type() : glsl_uint64_t_type());
749 else
750 val->type->type = (signedness ? glsl_int_type() : glsl_uint_type());
751 break;
752 }
753 case SpvOpTypeFloat: {
754 int bit_size = w[2];
755 val->type->base_type = vtn_base_type_scalar;
756 val->type->type = bit_size == 64 ? glsl_double_type() : glsl_float_type();
757 break;
758 }
759
760 case SpvOpTypeVector: {
761 struct vtn_type *base = vtn_value(b, w[2], vtn_value_type_type)->type;
762 unsigned elems = w[3];
763
764 assert(glsl_type_is_scalar(base->type));
765 val->type->base_type = vtn_base_type_vector;
766 val->type->type = glsl_vector_type(glsl_get_base_type(base->type), elems);
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 SpvOpAtomicCompareExchange:
2017 case SpvOpAtomicIIncrement:
2018 case SpvOpAtomicIDecrement:
2019 case SpvOpAtomicExchange:
2020 case SpvOpAtomicIAdd:
2021 case SpvOpAtomicSMin:
2022 case SpvOpAtomicUMin:
2023 case SpvOpAtomicSMax:
2024 case SpvOpAtomicUMax:
2025 case SpvOpAtomicAnd:
2026 case SpvOpAtomicOr:
2027 case SpvOpAtomicXor:
2028 fill_common_atomic_sources(b, opcode, w, &intrin->src[2]);
2029 break;
2030
2031 default:
2032 unreachable("Invalid image opcode");
2033 }
2034
2035 if (opcode != SpvOpImageWrite) {
2036 struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_ssa);
2037 struct vtn_type *type = vtn_value(b, w[1], vtn_value_type_type)->type;
2038
2039 unsigned dest_components =
2040 nir_intrinsic_infos[intrin->intrinsic].dest_components;
2041 if (intrin->intrinsic == nir_intrinsic_image_size) {
2042 dest_components = intrin->num_components =
2043 glsl_get_vector_elements(type->type);
2044 }
2045
2046 nir_ssa_dest_init(&intrin->instr, &intrin->dest,
2047 dest_components, 32, NULL);
2048
2049 nir_builder_instr_insert(&b->nb, &intrin->instr);
2050
2051 val->ssa = vtn_create_ssa_value(b, type->type);
2052 val->ssa->def = &intrin->dest.ssa;
2053 } else {
2054 nir_builder_instr_insert(&b->nb, &intrin->instr);
2055 }
2056 }
2057
2058 static nir_intrinsic_op
2059 get_ssbo_nir_atomic_op(SpvOp opcode)
2060 {
2061 switch (opcode) {
2062 case SpvOpAtomicLoad: return nir_intrinsic_load_ssbo;
2063 case SpvOpAtomicStore: return nir_intrinsic_store_ssbo;
2064 #define OP(S, N) case SpvOp##S: return nir_intrinsic_ssbo_##N;
2065 OP(AtomicExchange, atomic_exchange)
2066 OP(AtomicCompareExchange, atomic_comp_swap)
2067 OP(AtomicIIncrement, atomic_add)
2068 OP(AtomicIDecrement, atomic_add)
2069 OP(AtomicIAdd, atomic_add)
2070 OP(AtomicISub, atomic_add)
2071 OP(AtomicSMin, atomic_imin)
2072 OP(AtomicUMin, atomic_umin)
2073 OP(AtomicSMax, atomic_imax)
2074 OP(AtomicUMax, atomic_umax)
2075 OP(AtomicAnd, atomic_and)
2076 OP(AtomicOr, atomic_or)
2077 OP(AtomicXor, atomic_xor)
2078 #undef OP
2079 default:
2080 unreachable("Invalid SSBO atomic");
2081 }
2082 }
2083
2084 static nir_intrinsic_op
2085 get_shared_nir_atomic_op(SpvOp opcode)
2086 {
2087 switch (opcode) {
2088 case SpvOpAtomicLoad: return nir_intrinsic_load_var;
2089 case SpvOpAtomicStore: return nir_intrinsic_store_var;
2090 #define OP(S, N) case SpvOp##S: return nir_intrinsic_var_##N;
2091 OP(AtomicExchange, atomic_exchange)
2092 OP(AtomicCompareExchange, atomic_comp_swap)
2093 OP(AtomicIIncrement, atomic_add)
2094 OP(AtomicIDecrement, atomic_add)
2095 OP(AtomicIAdd, atomic_add)
2096 OP(AtomicISub, atomic_add)
2097 OP(AtomicSMin, atomic_imin)
2098 OP(AtomicUMin, atomic_umin)
2099 OP(AtomicSMax, atomic_imax)
2100 OP(AtomicUMax, atomic_umax)
2101 OP(AtomicAnd, atomic_and)
2102 OP(AtomicOr, atomic_or)
2103 OP(AtomicXor, atomic_xor)
2104 #undef OP
2105 default:
2106 unreachable("Invalid shared atomic");
2107 }
2108 }
2109
2110 static void
2111 vtn_handle_ssbo_or_shared_atomic(struct vtn_builder *b, SpvOp opcode,
2112 const uint32_t *w, unsigned count)
2113 {
2114 struct vtn_pointer *ptr;
2115 nir_intrinsic_instr *atomic;
2116
2117 switch (opcode) {
2118 case SpvOpAtomicLoad:
2119 case SpvOpAtomicExchange:
2120 case SpvOpAtomicCompareExchange:
2121 case SpvOpAtomicCompareExchangeWeak:
2122 case SpvOpAtomicIIncrement:
2123 case SpvOpAtomicIDecrement:
2124 case SpvOpAtomicIAdd:
2125 case SpvOpAtomicISub:
2126 case SpvOpAtomicSMin:
2127 case SpvOpAtomicUMin:
2128 case SpvOpAtomicSMax:
2129 case SpvOpAtomicUMax:
2130 case SpvOpAtomicAnd:
2131 case SpvOpAtomicOr:
2132 case SpvOpAtomicXor:
2133 ptr = vtn_value(b, w[3], vtn_value_type_pointer)->pointer;
2134 break;
2135
2136 case SpvOpAtomicStore:
2137 ptr = vtn_value(b, w[1], vtn_value_type_pointer)->pointer;
2138 break;
2139
2140 default:
2141 unreachable("Invalid SPIR-V atomic");
2142 }
2143
2144 /*
2145 SpvScope scope = w[4];
2146 SpvMemorySemanticsMask semantics = w[5];
2147 */
2148
2149 if (ptr->mode == vtn_variable_mode_workgroup) {
2150 nir_deref_var *deref = vtn_pointer_to_deref(b, ptr);
2151 const struct glsl_type *deref_type = nir_deref_tail(&deref->deref)->type;
2152 nir_intrinsic_op op = get_shared_nir_atomic_op(opcode);
2153 atomic = nir_intrinsic_instr_create(b->nb.shader, op);
2154 atomic->variables[0] = nir_deref_var_clone(deref, atomic);
2155
2156 switch (opcode) {
2157 case SpvOpAtomicLoad:
2158 atomic->num_components = glsl_get_vector_elements(deref_type);
2159 break;
2160
2161 case SpvOpAtomicStore:
2162 atomic->num_components = glsl_get_vector_elements(deref_type);
2163 nir_intrinsic_set_write_mask(atomic, (1 << atomic->num_components) - 1);
2164 atomic->src[0] = nir_src_for_ssa(vtn_ssa_value(b, w[4])->def);
2165 break;
2166
2167 case SpvOpAtomicExchange:
2168 case SpvOpAtomicCompareExchange:
2169 case SpvOpAtomicCompareExchangeWeak:
2170 case SpvOpAtomicIIncrement:
2171 case SpvOpAtomicIDecrement:
2172 case SpvOpAtomicIAdd:
2173 case SpvOpAtomicISub:
2174 case SpvOpAtomicSMin:
2175 case SpvOpAtomicUMin:
2176 case SpvOpAtomicSMax:
2177 case SpvOpAtomicUMax:
2178 case SpvOpAtomicAnd:
2179 case SpvOpAtomicOr:
2180 case SpvOpAtomicXor:
2181 fill_common_atomic_sources(b, opcode, w, &atomic->src[0]);
2182 break;
2183
2184 default:
2185 unreachable("Invalid SPIR-V atomic");
2186
2187 }
2188 } else {
2189 assert(ptr->mode == vtn_variable_mode_ssbo);
2190 nir_ssa_def *offset, *index;
2191 offset = vtn_pointer_to_offset(b, ptr, &index, NULL);
2192
2193 nir_intrinsic_op op = get_ssbo_nir_atomic_op(opcode);
2194
2195 atomic = nir_intrinsic_instr_create(b->nb.shader, op);
2196
2197 switch (opcode) {
2198 case SpvOpAtomicLoad:
2199 atomic->num_components = glsl_get_vector_elements(ptr->type->type);
2200 atomic->src[0] = nir_src_for_ssa(index);
2201 atomic->src[1] = nir_src_for_ssa(offset);
2202 break;
2203
2204 case SpvOpAtomicStore:
2205 atomic->num_components = glsl_get_vector_elements(ptr->type->type);
2206 nir_intrinsic_set_write_mask(atomic, (1 << atomic->num_components) - 1);
2207 atomic->src[0] = nir_src_for_ssa(vtn_ssa_value(b, w[4])->def);
2208 atomic->src[1] = nir_src_for_ssa(index);
2209 atomic->src[2] = nir_src_for_ssa(offset);
2210 break;
2211
2212 case SpvOpAtomicExchange:
2213 case SpvOpAtomicCompareExchange:
2214 case SpvOpAtomicCompareExchangeWeak:
2215 case SpvOpAtomicIIncrement:
2216 case SpvOpAtomicIDecrement:
2217 case SpvOpAtomicIAdd:
2218 case SpvOpAtomicISub:
2219 case SpvOpAtomicSMin:
2220 case SpvOpAtomicUMin:
2221 case SpvOpAtomicSMax:
2222 case SpvOpAtomicUMax:
2223 case SpvOpAtomicAnd:
2224 case SpvOpAtomicOr:
2225 case SpvOpAtomicXor:
2226 atomic->src[0] = nir_src_for_ssa(index);
2227 atomic->src[1] = nir_src_for_ssa(offset);
2228 fill_common_atomic_sources(b, opcode, w, &atomic->src[2]);
2229 break;
2230
2231 default:
2232 unreachable("Invalid SPIR-V atomic");
2233 }
2234 }
2235
2236 if (opcode != SpvOpAtomicStore) {
2237 struct vtn_type *type = vtn_value(b, w[1], vtn_value_type_type)->type;
2238
2239 nir_ssa_dest_init(&atomic->instr, &atomic->dest,
2240 glsl_get_vector_elements(type->type),
2241 glsl_get_bit_size(type->type), NULL);
2242
2243 struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_ssa);
2244 val->ssa = rzalloc(b, struct vtn_ssa_value);
2245 val->ssa->def = &atomic->dest.ssa;
2246 val->ssa->type = type->type;
2247 }
2248
2249 nir_builder_instr_insert(&b->nb, &atomic->instr);
2250 }
2251
2252 static nir_alu_instr *
2253 create_vec(nir_shader *shader, unsigned num_components, unsigned bit_size)
2254 {
2255 nir_op op;
2256 switch (num_components) {
2257 case 1: op = nir_op_fmov; break;
2258 case 2: op = nir_op_vec2; break;
2259 case 3: op = nir_op_vec3; break;
2260 case 4: op = nir_op_vec4; break;
2261 default: unreachable("bad vector size");
2262 }
2263
2264 nir_alu_instr *vec = nir_alu_instr_create(shader, op);
2265 nir_ssa_dest_init(&vec->instr, &vec->dest.dest, num_components,
2266 bit_size, NULL);
2267 vec->dest.write_mask = (1 << num_components) - 1;
2268
2269 return vec;
2270 }
2271
2272 struct vtn_ssa_value *
2273 vtn_ssa_transpose(struct vtn_builder *b, struct vtn_ssa_value *src)
2274 {
2275 if (src->transposed)
2276 return src->transposed;
2277
2278 struct vtn_ssa_value *dest =
2279 vtn_create_ssa_value(b, glsl_transposed_type(src->type));
2280
2281 for (unsigned i = 0; i < glsl_get_matrix_columns(dest->type); i++) {
2282 nir_alu_instr *vec = create_vec(b->shader,
2283 glsl_get_matrix_columns(src->type),
2284 glsl_get_bit_size(src->type));
2285 if (glsl_type_is_vector_or_scalar(src->type)) {
2286 vec->src[0].src = nir_src_for_ssa(src->def);
2287 vec->src[0].swizzle[0] = i;
2288 } else {
2289 for (unsigned j = 0; j < glsl_get_matrix_columns(src->type); j++) {
2290 vec->src[j].src = nir_src_for_ssa(src->elems[j]->def);
2291 vec->src[j].swizzle[0] = i;
2292 }
2293 }
2294 nir_builder_instr_insert(&b->nb, &vec->instr);
2295 dest->elems[i]->def = &vec->dest.dest.ssa;
2296 }
2297
2298 dest->transposed = src;
2299
2300 return dest;
2301 }
2302
2303 nir_ssa_def *
2304 vtn_vector_extract(struct vtn_builder *b, nir_ssa_def *src, unsigned index)
2305 {
2306 unsigned swiz[4] = { index };
2307 return nir_swizzle(&b->nb, src, swiz, 1, true);
2308 }
2309
2310 nir_ssa_def *
2311 vtn_vector_insert(struct vtn_builder *b, nir_ssa_def *src, nir_ssa_def *insert,
2312 unsigned index)
2313 {
2314 nir_alu_instr *vec = create_vec(b->shader, src->num_components,
2315 src->bit_size);
2316
2317 for (unsigned i = 0; i < src->num_components; i++) {
2318 if (i == index) {
2319 vec->src[i].src = nir_src_for_ssa(insert);
2320 } else {
2321 vec->src[i].src = nir_src_for_ssa(src);
2322 vec->src[i].swizzle[0] = i;
2323 }
2324 }
2325
2326 nir_builder_instr_insert(&b->nb, &vec->instr);
2327
2328 return &vec->dest.dest.ssa;
2329 }
2330
2331 nir_ssa_def *
2332 vtn_vector_extract_dynamic(struct vtn_builder *b, nir_ssa_def *src,
2333 nir_ssa_def *index)
2334 {
2335 nir_ssa_def *dest = vtn_vector_extract(b, src, 0);
2336 for (unsigned i = 1; i < src->num_components; i++)
2337 dest = nir_bcsel(&b->nb, nir_ieq(&b->nb, index, nir_imm_int(&b->nb, i)),
2338 vtn_vector_extract(b, src, i), dest);
2339
2340 return dest;
2341 }
2342
2343 nir_ssa_def *
2344 vtn_vector_insert_dynamic(struct vtn_builder *b, nir_ssa_def *src,
2345 nir_ssa_def *insert, nir_ssa_def *index)
2346 {
2347 nir_ssa_def *dest = vtn_vector_insert(b, src, insert, 0);
2348 for (unsigned i = 1; i < src->num_components; i++)
2349 dest = nir_bcsel(&b->nb, nir_ieq(&b->nb, index, nir_imm_int(&b->nb, i)),
2350 vtn_vector_insert(b, src, insert, i), dest);
2351
2352 return dest;
2353 }
2354
2355 static nir_ssa_def *
2356 vtn_vector_shuffle(struct vtn_builder *b, unsigned num_components,
2357 nir_ssa_def *src0, nir_ssa_def *src1,
2358 const uint32_t *indices)
2359 {
2360 nir_alu_instr *vec = create_vec(b->shader, num_components, src0->bit_size);
2361
2362 for (unsigned i = 0; i < num_components; i++) {
2363 uint32_t index = indices[i];
2364 if (index == 0xffffffff) {
2365 vec->src[i].src =
2366 nir_src_for_ssa(nir_ssa_undef(&b->nb, 1, src0->bit_size));
2367 } else if (index < src0->num_components) {
2368 vec->src[i].src = nir_src_for_ssa(src0);
2369 vec->src[i].swizzle[0] = index;
2370 } else {
2371 vec->src[i].src = nir_src_for_ssa(src1);
2372 vec->src[i].swizzle[0] = index - src0->num_components;
2373 }
2374 }
2375
2376 nir_builder_instr_insert(&b->nb, &vec->instr);
2377
2378 return &vec->dest.dest.ssa;
2379 }
2380
2381 /*
2382 * Concatentates a number of vectors/scalars together to produce a vector
2383 */
2384 static nir_ssa_def *
2385 vtn_vector_construct(struct vtn_builder *b, unsigned num_components,
2386 unsigned num_srcs, nir_ssa_def **srcs)
2387 {
2388 nir_alu_instr *vec = create_vec(b->shader, num_components,
2389 srcs[0]->bit_size);
2390
2391 /* From the SPIR-V 1.1 spec for OpCompositeConstruct:
2392 *
2393 * "When constructing a vector, there must be at least two Constituent
2394 * operands."
2395 */
2396 assert(num_srcs >= 2);
2397
2398 unsigned dest_idx = 0;
2399 for (unsigned i = 0; i < num_srcs; i++) {
2400 nir_ssa_def *src = srcs[i];
2401 assert(dest_idx + src->num_components <= num_components);
2402 for (unsigned j = 0; j < src->num_components; j++) {
2403 vec->src[dest_idx].src = nir_src_for_ssa(src);
2404 vec->src[dest_idx].swizzle[0] = j;
2405 dest_idx++;
2406 }
2407 }
2408
2409 /* From the SPIR-V 1.1 spec for OpCompositeConstruct:
2410 *
2411 * "When constructing a vector, the total number of components in all
2412 * the operands must equal the number of components in Result Type."
2413 */
2414 assert(dest_idx == num_components);
2415
2416 nir_builder_instr_insert(&b->nb, &vec->instr);
2417
2418 return &vec->dest.dest.ssa;
2419 }
2420
2421 static struct vtn_ssa_value *
2422 vtn_composite_copy(void *mem_ctx, struct vtn_ssa_value *src)
2423 {
2424 struct vtn_ssa_value *dest = rzalloc(mem_ctx, struct vtn_ssa_value);
2425 dest->type = src->type;
2426
2427 if (glsl_type_is_vector_or_scalar(src->type)) {
2428 dest->def = src->def;
2429 } else {
2430 unsigned elems = glsl_get_length(src->type);
2431
2432 dest->elems = ralloc_array(mem_ctx, struct vtn_ssa_value *, elems);
2433 for (unsigned i = 0; i < elems; i++)
2434 dest->elems[i] = vtn_composite_copy(mem_ctx, src->elems[i]);
2435 }
2436
2437 return dest;
2438 }
2439
2440 static struct vtn_ssa_value *
2441 vtn_composite_insert(struct vtn_builder *b, struct vtn_ssa_value *src,
2442 struct vtn_ssa_value *insert, const uint32_t *indices,
2443 unsigned num_indices)
2444 {
2445 struct vtn_ssa_value *dest = vtn_composite_copy(b, src);
2446
2447 struct vtn_ssa_value *cur = dest;
2448 unsigned i;
2449 for (i = 0; i < num_indices - 1; i++) {
2450 cur = cur->elems[indices[i]];
2451 }
2452
2453 if (glsl_type_is_vector_or_scalar(cur->type)) {
2454 /* According to the SPIR-V spec, OpCompositeInsert may work down to
2455 * the component granularity. In that case, the last index will be
2456 * the index to insert the scalar into the vector.
2457 */
2458
2459 cur->def = vtn_vector_insert(b, cur->def, insert->def, indices[i]);
2460 } else {
2461 cur->elems[indices[i]] = insert;
2462 }
2463
2464 return dest;
2465 }
2466
2467 static struct vtn_ssa_value *
2468 vtn_composite_extract(struct vtn_builder *b, struct vtn_ssa_value *src,
2469 const uint32_t *indices, unsigned num_indices)
2470 {
2471 struct vtn_ssa_value *cur = src;
2472 for (unsigned i = 0; i < num_indices; i++) {
2473 if (glsl_type_is_vector_or_scalar(cur->type)) {
2474 assert(i == num_indices - 1);
2475 /* According to the SPIR-V spec, OpCompositeExtract may work down to
2476 * the component granularity. The last index will be the index of the
2477 * vector to extract.
2478 */
2479
2480 struct vtn_ssa_value *ret = rzalloc(b, struct vtn_ssa_value);
2481 ret->type = glsl_scalar_type(glsl_get_base_type(cur->type));
2482 ret->def = vtn_vector_extract(b, cur->def, indices[i]);
2483 return ret;
2484 } else {
2485 cur = cur->elems[indices[i]];
2486 }
2487 }
2488
2489 return cur;
2490 }
2491
2492 static void
2493 vtn_handle_composite(struct vtn_builder *b, SpvOp opcode,
2494 const uint32_t *w, unsigned count)
2495 {
2496 struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_ssa);
2497 const struct glsl_type *type =
2498 vtn_value(b, w[1], vtn_value_type_type)->type->type;
2499 val->ssa = vtn_create_ssa_value(b, type);
2500
2501 switch (opcode) {
2502 case SpvOpVectorExtractDynamic:
2503 val->ssa->def = vtn_vector_extract_dynamic(b, vtn_ssa_value(b, w[3])->def,
2504 vtn_ssa_value(b, w[4])->def);
2505 break;
2506
2507 case SpvOpVectorInsertDynamic:
2508 val->ssa->def = vtn_vector_insert_dynamic(b, vtn_ssa_value(b, w[3])->def,
2509 vtn_ssa_value(b, w[4])->def,
2510 vtn_ssa_value(b, w[5])->def);
2511 break;
2512
2513 case SpvOpVectorShuffle:
2514 val->ssa->def = vtn_vector_shuffle(b, glsl_get_vector_elements(type),
2515 vtn_ssa_value(b, w[3])->def,
2516 vtn_ssa_value(b, w[4])->def,
2517 w + 5);
2518 break;
2519
2520 case SpvOpCompositeConstruct: {
2521 unsigned elems = count - 3;
2522 if (glsl_type_is_vector_or_scalar(type)) {
2523 nir_ssa_def *srcs[4];
2524 for (unsigned i = 0; i < elems; i++)
2525 srcs[i] = vtn_ssa_value(b, w[3 + i])->def;
2526 val->ssa->def =
2527 vtn_vector_construct(b, glsl_get_vector_elements(type),
2528 elems, srcs);
2529 } else {
2530 val->ssa->elems = ralloc_array(b, struct vtn_ssa_value *, elems);
2531 for (unsigned i = 0; i < elems; i++)
2532 val->ssa->elems[i] = vtn_ssa_value(b, w[3 + i]);
2533 }
2534 break;
2535 }
2536 case SpvOpCompositeExtract:
2537 val->ssa = vtn_composite_extract(b, vtn_ssa_value(b, w[3]),
2538 w + 4, count - 4);
2539 break;
2540
2541 case SpvOpCompositeInsert:
2542 val->ssa = vtn_composite_insert(b, vtn_ssa_value(b, w[4]),
2543 vtn_ssa_value(b, w[3]),
2544 w + 5, count - 5);
2545 break;
2546
2547 case SpvOpCopyObject:
2548 val->ssa = vtn_composite_copy(b, vtn_ssa_value(b, w[3]));
2549 break;
2550
2551 default:
2552 unreachable("unknown composite operation");
2553 }
2554 }
2555
2556 static void
2557 vtn_handle_barrier(struct vtn_builder *b, SpvOp opcode,
2558 const uint32_t *w, unsigned count)
2559 {
2560 nir_intrinsic_op intrinsic_op;
2561 switch (opcode) {
2562 case SpvOpEmitVertex:
2563 case SpvOpEmitStreamVertex:
2564 intrinsic_op = nir_intrinsic_emit_vertex;
2565 break;
2566 case SpvOpEndPrimitive:
2567 case SpvOpEndStreamPrimitive:
2568 intrinsic_op = nir_intrinsic_end_primitive;
2569 break;
2570 case SpvOpMemoryBarrier:
2571 intrinsic_op = nir_intrinsic_memory_barrier;
2572 break;
2573 case SpvOpControlBarrier:
2574 intrinsic_op = nir_intrinsic_barrier;
2575 break;
2576 default:
2577 unreachable("unknown barrier instruction");
2578 }
2579
2580 nir_intrinsic_instr *intrin =
2581 nir_intrinsic_instr_create(b->shader, intrinsic_op);
2582
2583 if (opcode == SpvOpEmitStreamVertex || opcode == SpvOpEndStreamPrimitive)
2584 nir_intrinsic_set_stream_id(intrin, w[1]);
2585
2586 nir_builder_instr_insert(&b->nb, &intrin->instr);
2587 }
2588
2589 static unsigned
2590 gl_primitive_from_spv_execution_mode(SpvExecutionMode mode)
2591 {
2592 switch (mode) {
2593 case SpvExecutionModeInputPoints:
2594 case SpvExecutionModeOutputPoints:
2595 return 0; /* GL_POINTS */
2596 case SpvExecutionModeInputLines:
2597 return 1; /* GL_LINES */
2598 case SpvExecutionModeInputLinesAdjacency:
2599 return 0x000A; /* GL_LINE_STRIP_ADJACENCY_ARB */
2600 case SpvExecutionModeTriangles:
2601 return 4; /* GL_TRIANGLES */
2602 case SpvExecutionModeInputTrianglesAdjacency:
2603 return 0x000C; /* GL_TRIANGLES_ADJACENCY_ARB */
2604 case SpvExecutionModeQuads:
2605 return 7; /* GL_QUADS */
2606 case SpvExecutionModeIsolines:
2607 return 0x8E7A; /* GL_ISOLINES */
2608 case SpvExecutionModeOutputLineStrip:
2609 return 3; /* GL_LINE_STRIP */
2610 case SpvExecutionModeOutputTriangleStrip:
2611 return 5; /* GL_TRIANGLE_STRIP */
2612 default:
2613 assert(!"Invalid primitive type");
2614 return 4;
2615 }
2616 }
2617
2618 static unsigned
2619 vertices_in_from_spv_execution_mode(SpvExecutionMode mode)
2620 {
2621 switch (mode) {
2622 case SpvExecutionModeInputPoints:
2623 return 1;
2624 case SpvExecutionModeInputLines:
2625 return 2;
2626 case SpvExecutionModeInputLinesAdjacency:
2627 return 4;
2628 case SpvExecutionModeTriangles:
2629 return 3;
2630 case SpvExecutionModeInputTrianglesAdjacency:
2631 return 6;
2632 default:
2633 assert(!"Invalid GS input mode");
2634 return 0;
2635 }
2636 }
2637
2638 static gl_shader_stage
2639 stage_for_execution_model(SpvExecutionModel model)
2640 {
2641 switch (model) {
2642 case SpvExecutionModelVertex:
2643 return MESA_SHADER_VERTEX;
2644 case SpvExecutionModelTessellationControl:
2645 return MESA_SHADER_TESS_CTRL;
2646 case SpvExecutionModelTessellationEvaluation:
2647 return MESA_SHADER_TESS_EVAL;
2648 case SpvExecutionModelGeometry:
2649 return MESA_SHADER_GEOMETRY;
2650 case SpvExecutionModelFragment:
2651 return MESA_SHADER_FRAGMENT;
2652 case SpvExecutionModelGLCompute:
2653 return MESA_SHADER_COMPUTE;
2654 default:
2655 unreachable("Unsupported execution model");
2656 }
2657 }
2658
2659 #define spv_check_supported(name, cap) do { \
2660 if (!(b->ext && b->ext->name)) \
2661 vtn_warn("Unsupported SPIR-V capability: %s", \
2662 spirv_capability_to_string(cap)); \
2663 } while(0)
2664
2665 static bool
2666 vtn_handle_preamble_instruction(struct vtn_builder *b, SpvOp opcode,
2667 const uint32_t *w, unsigned count)
2668 {
2669 switch (opcode) {
2670 case SpvOpSource:
2671 case SpvOpSourceExtension:
2672 case SpvOpSourceContinued:
2673 case SpvOpExtension:
2674 /* Unhandled, but these are for debug so that's ok. */
2675 break;
2676
2677 case SpvOpCapability: {
2678 SpvCapability cap = w[1];
2679 switch (cap) {
2680 case SpvCapabilityMatrix:
2681 case SpvCapabilityShader:
2682 case SpvCapabilityGeometry:
2683 case SpvCapabilityGeometryPointSize:
2684 case SpvCapabilityUniformBufferArrayDynamicIndexing:
2685 case SpvCapabilitySampledImageArrayDynamicIndexing:
2686 case SpvCapabilityStorageBufferArrayDynamicIndexing:
2687 case SpvCapabilityStorageImageArrayDynamicIndexing:
2688 case SpvCapabilityImageRect:
2689 case SpvCapabilitySampledRect:
2690 case SpvCapabilitySampled1D:
2691 case SpvCapabilityImage1D:
2692 case SpvCapabilitySampledCubeArray:
2693 case SpvCapabilitySampledBuffer:
2694 case SpvCapabilityImageBuffer:
2695 case SpvCapabilityImageQuery:
2696 case SpvCapabilityDerivativeControl:
2697 case SpvCapabilityInterpolationFunction:
2698 case SpvCapabilityMultiViewport:
2699 case SpvCapabilitySampleRateShading:
2700 case SpvCapabilityClipDistance:
2701 case SpvCapabilityCullDistance:
2702 case SpvCapabilityInputAttachment:
2703 case SpvCapabilityImageGatherExtended:
2704 case SpvCapabilityStorageImageExtendedFormats:
2705 break;
2706
2707 case SpvCapabilityGeometryStreams:
2708 case SpvCapabilityLinkage:
2709 case SpvCapabilityVector16:
2710 case SpvCapabilityFloat16Buffer:
2711 case SpvCapabilityFloat16:
2712 case SpvCapabilityInt64Atomics:
2713 case SpvCapabilityAtomicStorage:
2714 case SpvCapabilityInt16:
2715 case SpvCapabilityStorageImageMultisample:
2716 case SpvCapabilityImageCubeArray:
2717 case SpvCapabilityInt8:
2718 case SpvCapabilitySparseResidency:
2719 case SpvCapabilityMinLod:
2720 case SpvCapabilityTransformFeedback:
2721 vtn_warn("Unsupported SPIR-V capability: %s",
2722 spirv_capability_to_string(cap));
2723 break;
2724
2725 case SpvCapabilityFloat64:
2726 spv_check_supported(float64, cap);
2727 break;
2728 case SpvCapabilityInt64:
2729 spv_check_supported(int64, cap);
2730 break;
2731
2732 case SpvCapabilityAddresses:
2733 case SpvCapabilityKernel:
2734 case SpvCapabilityImageBasic:
2735 case SpvCapabilityImageReadWrite:
2736 case SpvCapabilityImageMipmap:
2737 case SpvCapabilityPipes:
2738 case SpvCapabilityGroups:
2739 case SpvCapabilityDeviceEnqueue:
2740 case SpvCapabilityLiteralSampler:
2741 case SpvCapabilityGenericPointer:
2742 vtn_warn("Unsupported OpenCL-style SPIR-V capability: %s",
2743 spirv_capability_to_string(cap));
2744 break;
2745
2746 case SpvCapabilityImageMSArray:
2747 spv_check_supported(image_ms_array, cap);
2748 break;
2749
2750 case SpvCapabilityTessellation:
2751 case SpvCapabilityTessellationPointSize:
2752 spv_check_supported(tessellation, cap);
2753 break;
2754
2755 case SpvCapabilityDrawParameters:
2756 spv_check_supported(draw_parameters, cap);
2757 break;
2758
2759 case SpvCapabilityStorageImageReadWithoutFormat:
2760 spv_check_supported(image_read_without_format, cap);
2761 break;
2762
2763 case SpvCapabilityStorageImageWriteWithoutFormat:
2764 spv_check_supported(image_write_without_format, cap);
2765 break;
2766
2767 case SpvCapabilityMultiView:
2768 spv_check_supported(multiview, cap);
2769 break;
2770
2771 default:
2772 unreachable("Unhandled capability");
2773 }
2774 break;
2775 }
2776
2777 case SpvOpExtInstImport:
2778 vtn_handle_extension(b, opcode, w, count);
2779 break;
2780
2781 case SpvOpMemoryModel:
2782 assert(w[1] == SpvAddressingModelLogical);
2783 assert(w[2] == SpvMemoryModelGLSL450);
2784 break;
2785
2786 case SpvOpEntryPoint: {
2787 struct vtn_value *entry_point = &b->values[w[2]];
2788 /* Let this be a name label regardless */
2789 unsigned name_words;
2790 entry_point->name = vtn_string_literal(b, &w[3], count - 3, &name_words);
2791
2792 if (strcmp(entry_point->name, b->entry_point_name) != 0 ||
2793 stage_for_execution_model(w[1]) != b->entry_point_stage)
2794 break;
2795
2796 assert(b->entry_point == NULL);
2797 b->entry_point = entry_point;
2798 break;
2799 }
2800
2801 case SpvOpString:
2802 vtn_push_value(b, w[1], vtn_value_type_string)->str =
2803 vtn_string_literal(b, &w[2], count - 2, NULL);
2804 break;
2805
2806 case SpvOpName:
2807 b->values[w[1]].name = vtn_string_literal(b, &w[2], count - 2, NULL);
2808 break;
2809
2810 case SpvOpMemberName:
2811 /* TODO */
2812 break;
2813
2814 case SpvOpExecutionMode:
2815 case SpvOpDecorationGroup:
2816 case SpvOpDecorate:
2817 case SpvOpMemberDecorate:
2818 case SpvOpGroupDecorate:
2819 case SpvOpGroupMemberDecorate:
2820 vtn_handle_decoration(b, opcode, w, count);
2821 break;
2822
2823 default:
2824 return false; /* End of preamble */
2825 }
2826
2827 return true;
2828 }
2829
2830 static void
2831 vtn_handle_execution_mode(struct vtn_builder *b, struct vtn_value *entry_point,
2832 const struct vtn_decoration *mode, void *data)
2833 {
2834 assert(b->entry_point == entry_point);
2835
2836 switch(mode->exec_mode) {
2837 case SpvExecutionModeOriginUpperLeft:
2838 case SpvExecutionModeOriginLowerLeft:
2839 b->origin_upper_left =
2840 (mode->exec_mode == SpvExecutionModeOriginUpperLeft);
2841 break;
2842
2843 case SpvExecutionModeEarlyFragmentTests:
2844 assert(b->shader->stage == MESA_SHADER_FRAGMENT);
2845 b->shader->info.fs.early_fragment_tests = true;
2846 break;
2847
2848 case SpvExecutionModeInvocations:
2849 assert(b->shader->stage == MESA_SHADER_GEOMETRY);
2850 b->shader->info.gs.invocations = MAX2(1, mode->literals[0]);
2851 break;
2852
2853 case SpvExecutionModeDepthReplacing:
2854 assert(b->shader->stage == MESA_SHADER_FRAGMENT);
2855 b->shader->info.fs.depth_layout = FRAG_DEPTH_LAYOUT_ANY;
2856 break;
2857 case SpvExecutionModeDepthGreater:
2858 assert(b->shader->stage == MESA_SHADER_FRAGMENT);
2859 b->shader->info.fs.depth_layout = FRAG_DEPTH_LAYOUT_GREATER;
2860 break;
2861 case SpvExecutionModeDepthLess:
2862 assert(b->shader->stage == MESA_SHADER_FRAGMENT);
2863 b->shader->info.fs.depth_layout = FRAG_DEPTH_LAYOUT_LESS;
2864 break;
2865 case SpvExecutionModeDepthUnchanged:
2866 assert(b->shader->stage == MESA_SHADER_FRAGMENT);
2867 b->shader->info.fs.depth_layout = FRAG_DEPTH_LAYOUT_UNCHANGED;
2868 break;
2869
2870 case SpvExecutionModeLocalSize:
2871 assert(b->shader->stage == MESA_SHADER_COMPUTE);
2872 b->shader->info.cs.local_size[0] = mode->literals[0];
2873 b->shader->info.cs.local_size[1] = mode->literals[1];
2874 b->shader->info.cs.local_size[2] = mode->literals[2];
2875 break;
2876 case SpvExecutionModeLocalSizeHint:
2877 break; /* Nothing to do with this */
2878
2879 case SpvExecutionModeOutputVertices:
2880 if (b->shader->stage == MESA_SHADER_TESS_CTRL ||
2881 b->shader->stage == MESA_SHADER_TESS_EVAL) {
2882 b->shader->info.tess.tcs_vertices_out = mode->literals[0];
2883 } else {
2884 assert(b->shader->stage == MESA_SHADER_GEOMETRY);
2885 b->shader->info.gs.vertices_out = mode->literals[0];
2886 }
2887 break;
2888
2889 case SpvExecutionModeInputPoints:
2890 case SpvExecutionModeInputLines:
2891 case SpvExecutionModeInputLinesAdjacency:
2892 case SpvExecutionModeTriangles:
2893 case SpvExecutionModeInputTrianglesAdjacency:
2894 case SpvExecutionModeQuads:
2895 case SpvExecutionModeIsolines:
2896 if (b->shader->stage == MESA_SHADER_TESS_CTRL ||
2897 b->shader->stage == MESA_SHADER_TESS_EVAL) {
2898 b->shader->info.tess.primitive_mode =
2899 gl_primitive_from_spv_execution_mode(mode->exec_mode);
2900 } else {
2901 assert(b->shader->stage == MESA_SHADER_GEOMETRY);
2902 b->shader->info.gs.vertices_in =
2903 vertices_in_from_spv_execution_mode(mode->exec_mode);
2904 }
2905 break;
2906
2907 case SpvExecutionModeOutputPoints:
2908 case SpvExecutionModeOutputLineStrip:
2909 case SpvExecutionModeOutputTriangleStrip:
2910 assert(b->shader->stage == MESA_SHADER_GEOMETRY);
2911 b->shader->info.gs.output_primitive =
2912 gl_primitive_from_spv_execution_mode(mode->exec_mode);
2913 break;
2914
2915 case SpvExecutionModeSpacingEqual:
2916 assert(b->shader->stage == MESA_SHADER_TESS_CTRL ||
2917 b->shader->stage == MESA_SHADER_TESS_EVAL);
2918 b->shader->info.tess.spacing = TESS_SPACING_EQUAL;
2919 break;
2920 case SpvExecutionModeSpacingFractionalEven:
2921 assert(b->shader->stage == MESA_SHADER_TESS_CTRL ||
2922 b->shader->stage == MESA_SHADER_TESS_EVAL);
2923 b->shader->info.tess.spacing = TESS_SPACING_FRACTIONAL_EVEN;
2924 break;
2925 case SpvExecutionModeSpacingFractionalOdd:
2926 assert(b->shader->stage == MESA_SHADER_TESS_CTRL ||
2927 b->shader->stage == MESA_SHADER_TESS_EVAL);
2928 b->shader->info.tess.spacing = TESS_SPACING_FRACTIONAL_ODD;
2929 break;
2930 case SpvExecutionModeVertexOrderCw:
2931 assert(b->shader->stage == MESA_SHADER_TESS_CTRL ||
2932 b->shader->stage == MESA_SHADER_TESS_EVAL);
2933 /* Vulkan's notion of CCW seems to match the hardware backends,
2934 * but be the opposite of OpenGL. Currently NIR follows GL semantics,
2935 * so we set it backwards here.
2936 */
2937 b->shader->info.tess.ccw = true;
2938 break;
2939 case SpvExecutionModeVertexOrderCcw:
2940 assert(b->shader->stage == MESA_SHADER_TESS_CTRL ||
2941 b->shader->stage == MESA_SHADER_TESS_EVAL);
2942 /* Backwards; see above */
2943 b->shader->info.tess.ccw = false;
2944 break;
2945 case SpvExecutionModePointMode:
2946 assert(b->shader->stage == MESA_SHADER_TESS_CTRL ||
2947 b->shader->stage == MESA_SHADER_TESS_EVAL);
2948 b->shader->info.tess.point_mode = true;
2949 break;
2950
2951 case SpvExecutionModePixelCenterInteger:
2952 b->pixel_center_integer = true;
2953 break;
2954
2955 case SpvExecutionModeXfb:
2956 assert(!"Unhandled execution mode");
2957 break;
2958
2959 case SpvExecutionModeVecTypeHint:
2960 case SpvExecutionModeContractionOff:
2961 break; /* OpenCL */
2962
2963 default:
2964 unreachable("Unhandled execution mode");
2965 }
2966 }
2967
2968 static bool
2969 vtn_handle_variable_or_type_instruction(struct vtn_builder *b, SpvOp opcode,
2970 const uint32_t *w, unsigned count)
2971 {
2972 switch (opcode) {
2973 case SpvOpSource:
2974 case SpvOpSourceContinued:
2975 case SpvOpSourceExtension:
2976 case SpvOpExtension:
2977 case SpvOpCapability:
2978 case SpvOpExtInstImport:
2979 case SpvOpMemoryModel:
2980 case SpvOpEntryPoint:
2981 case SpvOpExecutionMode:
2982 case SpvOpString:
2983 case SpvOpName:
2984 case SpvOpMemberName:
2985 case SpvOpDecorationGroup:
2986 case SpvOpDecorate:
2987 case SpvOpMemberDecorate:
2988 case SpvOpGroupDecorate:
2989 case SpvOpGroupMemberDecorate:
2990 assert(!"Invalid opcode types and variables section");
2991 break;
2992
2993 case SpvOpTypeVoid:
2994 case SpvOpTypeBool:
2995 case SpvOpTypeInt:
2996 case SpvOpTypeFloat:
2997 case SpvOpTypeVector:
2998 case SpvOpTypeMatrix:
2999 case SpvOpTypeImage:
3000 case SpvOpTypeSampler:
3001 case SpvOpTypeSampledImage:
3002 case SpvOpTypeArray:
3003 case SpvOpTypeRuntimeArray:
3004 case SpvOpTypeStruct:
3005 case SpvOpTypeOpaque:
3006 case SpvOpTypePointer:
3007 case SpvOpTypeFunction:
3008 case SpvOpTypeEvent:
3009 case SpvOpTypeDeviceEvent:
3010 case SpvOpTypeReserveId:
3011 case SpvOpTypeQueue:
3012 case SpvOpTypePipe:
3013 vtn_handle_type(b, opcode, w, count);
3014 break;
3015
3016 case SpvOpConstantTrue:
3017 case SpvOpConstantFalse:
3018 case SpvOpConstant:
3019 case SpvOpConstantComposite:
3020 case SpvOpConstantSampler:
3021 case SpvOpConstantNull:
3022 case SpvOpSpecConstantTrue:
3023 case SpvOpSpecConstantFalse:
3024 case SpvOpSpecConstant:
3025 case SpvOpSpecConstantComposite:
3026 case SpvOpSpecConstantOp:
3027 vtn_handle_constant(b, opcode, w, count);
3028 break;
3029
3030 case SpvOpUndef:
3031 case SpvOpVariable:
3032 vtn_handle_variables(b, opcode, w, count);
3033 break;
3034
3035 default:
3036 return false; /* End of preamble */
3037 }
3038
3039 return true;
3040 }
3041
3042 static bool
3043 vtn_handle_body_instruction(struct vtn_builder *b, SpvOp opcode,
3044 const uint32_t *w, unsigned count)
3045 {
3046 switch (opcode) {
3047 case SpvOpLabel:
3048 break;
3049
3050 case SpvOpLoopMerge:
3051 case SpvOpSelectionMerge:
3052 /* This is handled by cfg pre-pass and walk_blocks */
3053 break;
3054
3055 case SpvOpUndef: {
3056 struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_undef);
3057 val->type = vtn_value(b, w[1], vtn_value_type_type)->type;
3058 break;
3059 }
3060
3061 case SpvOpExtInst:
3062 vtn_handle_extension(b, opcode, w, count);
3063 break;
3064
3065 case SpvOpVariable:
3066 case SpvOpLoad:
3067 case SpvOpStore:
3068 case SpvOpCopyMemory:
3069 case SpvOpCopyMemorySized:
3070 case SpvOpAccessChain:
3071 case SpvOpPtrAccessChain:
3072 case SpvOpInBoundsAccessChain:
3073 case SpvOpArrayLength:
3074 vtn_handle_variables(b, opcode, w, count);
3075 break;
3076
3077 case SpvOpFunctionCall:
3078 vtn_handle_function_call(b, opcode, w, count);
3079 break;
3080
3081 case SpvOpSampledImage:
3082 case SpvOpImage:
3083 case SpvOpImageSampleImplicitLod:
3084 case SpvOpImageSampleExplicitLod:
3085 case SpvOpImageSampleDrefImplicitLod:
3086 case SpvOpImageSampleDrefExplicitLod:
3087 case SpvOpImageSampleProjImplicitLod:
3088 case SpvOpImageSampleProjExplicitLod:
3089 case SpvOpImageSampleProjDrefImplicitLod:
3090 case SpvOpImageSampleProjDrefExplicitLod:
3091 case SpvOpImageFetch:
3092 case SpvOpImageGather:
3093 case SpvOpImageDrefGather:
3094 case SpvOpImageQuerySizeLod:
3095 case SpvOpImageQueryLod:
3096 case SpvOpImageQueryLevels:
3097 case SpvOpImageQuerySamples:
3098 vtn_handle_texture(b, opcode, w, count);
3099 break;
3100
3101 case SpvOpImageRead:
3102 case SpvOpImageWrite:
3103 case SpvOpImageTexelPointer:
3104 vtn_handle_image(b, opcode, w, count);
3105 break;
3106
3107 case SpvOpImageQuerySize: {
3108 struct vtn_pointer *image =
3109 vtn_value(b, w[3], vtn_value_type_pointer)->pointer;
3110 if (image->mode == vtn_variable_mode_image) {
3111 vtn_handle_image(b, opcode, w, count);
3112 } else {
3113 assert(image->mode == vtn_variable_mode_sampler);
3114 vtn_handle_texture(b, opcode, w, count);
3115 }
3116 break;
3117 }
3118
3119 case SpvOpAtomicLoad:
3120 case SpvOpAtomicExchange:
3121 case SpvOpAtomicCompareExchange:
3122 case SpvOpAtomicCompareExchangeWeak:
3123 case SpvOpAtomicIIncrement:
3124 case SpvOpAtomicIDecrement:
3125 case SpvOpAtomicIAdd:
3126 case SpvOpAtomicISub:
3127 case SpvOpAtomicSMin:
3128 case SpvOpAtomicUMin:
3129 case SpvOpAtomicSMax:
3130 case SpvOpAtomicUMax:
3131 case SpvOpAtomicAnd:
3132 case SpvOpAtomicOr:
3133 case SpvOpAtomicXor: {
3134 struct vtn_value *pointer = vtn_untyped_value(b, w[3]);
3135 if (pointer->value_type == vtn_value_type_image_pointer) {
3136 vtn_handle_image(b, opcode, w, count);
3137 } else {
3138 assert(pointer->value_type == vtn_value_type_pointer);
3139 vtn_handle_ssbo_or_shared_atomic(b, opcode, w, count);
3140 }
3141 break;
3142 }
3143
3144 case SpvOpAtomicStore: {
3145 struct vtn_value *pointer = vtn_untyped_value(b, w[1]);
3146 if (pointer->value_type == vtn_value_type_image_pointer) {
3147 vtn_handle_image(b, opcode, w, count);
3148 } else {
3149 assert(pointer->value_type == vtn_value_type_pointer);
3150 vtn_handle_ssbo_or_shared_atomic(b, opcode, w, count);
3151 }
3152 break;
3153 }
3154
3155 case SpvOpSNegate:
3156 case SpvOpFNegate:
3157 case SpvOpNot:
3158 case SpvOpAny:
3159 case SpvOpAll:
3160 case SpvOpConvertFToU:
3161 case SpvOpConvertFToS:
3162 case SpvOpConvertSToF:
3163 case SpvOpConvertUToF:
3164 case SpvOpUConvert:
3165 case SpvOpSConvert:
3166 case SpvOpFConvert:
3167 case SpvOpQuantizeToF16:
3168 case SpvOpConvertPtrToU:
3169 case SpvOpConvertUToPtr:
3170 case SpvOpPtrCastToGeneric:
3171 case SpvOpGenericCastToPtr:
3172 case SpvOpBitcast:
3173 case SpvOpIsNan:
3174 case SpvOpIsInf:
3175 case SpvOpIsFinite:
3176 case SpvOpIsNormal:
3177 case SpvOpSignBitSet:
3178 case SpvOpLessOrGreater:
3179 case SpvOpOrdered:
3180 case SpvOpUnordered:
3181 case SpvOpIAdd:
3182 case SpvOpFAdd:
3183 case SpvOpISub:
3184 case SpvOpFSub:
3185 case SpvOpIMul:
3186 case SpvOpFMul:
3187 case SpvOpUDiv:
3188 case SpvOpSDiv:
3189 case SpvOpFDiv:
3190 case SpvOpUMod:
3191 case SpvOpSRem:
3192 case SpvOpSMod:
3193 case SpvOpFRem:
3194 case SpvOpFMod:
3195 case SpvOpVectorTimesScalar:
3196 case SpvOpDot:
3197 case SpvOpIAddCarry:
3198 case SpvOpISubBorrow:
3199 case SpvOpUMulExtended:
3200 case SpvOpSMulExtended:
3201 case SpvOpShiftRightLogical:
3202 case SpvOpShiftRightArithmetic:
3203 case SpvOpShiftLeftLogical:
3204 case SpvOpLogicalEqual:
3205 case SpvOpLogicalNotEqual:
3206 case SpvOpLogicalOr:
3207 case SpvOpLogicalAnd:
3208 case SpvOpLogicalNot:
3209 case SpvOpBitwiseOr:
3210 case SpvOpBitwiseXor:
3211 case SpvOpBitwiseAnd:
3212 case SpvOpSelect:
3213 case SpvOpIEqual:
3214 case SpvOpFOrdEqual:
3215 case SpvOpFUnordEqual:
3216 case SpvOpINotEqual:
3217 case SpvOpFOrdNotEqual:
3218 case SpvOpFUnordNotEqual:
3219 case SpvOpULessThan:
3220 case SpvOpSLessThan:
3221 case SpvOpFOrdLessThan:
3222 case SpvOpFUnordLessThan:
3223 case SpvOpUGreaterThan:
3224 case SpvOpSGreaterThan:
3225 case SpvOpFOrdGreaterThan:
3226 case SpvOpFUnordGreaterThan:
3227 case SpvOpULessThanEqual:
3228 case SpvOpSLessThanEqual:
3229 case SpvOpFOrdLessThanEqual:
3230 case SpvOpFUnordLessThanEqual:
3231 case SpvOpUGreaterThanEqual:
3232 case SpvOpSGreaterThanEqual:
3233 case SpvOpFOrdGreaterThanEqual:
3234 case SpvOpFUnordGreaterThanEqual:
3235 case SpvOpDPdx:
3236 case SpvOpDPdy:
3237 case SpvOpFwidth:
3238 case SpvOpDPdxFine:
3239 case SpvOpDPdyFine:
3240 case SpvOpFwidthFine:
3241 case SpvOpDPdxCoarse:
3242 case SpvOpDPdyCoarse:
3243 case SpvOpFwidthCoarse:
3244 case SpvOpBitFieldInsert:
3245 case SpvOpBitFieldSExtract:
3246 case SpvOpBitFieldUExtract:
3247 case SpvOpBitReverse:
3248 case SpvOpBitCount:
3249 case SpvOpTranspose:
3250 case SpvOpOuterProduct:
3251 case SpvOpMatrixTimesScalar:
3252 case SpvOpVectorTimesMatrix:
3253 case SpvOpMatrixTimesVector:
3254 case SpvOpMatrixTimesMatrix:
3255 vtn_handle_alu(b, opcode, w, count);
3256 break;
3257
3258 case SpvOpVectorExtractDynamic:
3259 case SpvOpVectorInsertDynamic:
3260 case SpvOpVectorShuffle:
3261 case SpvOpCompositeConstruct:
3262 case SpvOpCompositeExtract:
3263 case SpvOpCompositeInsert:
3264 case SpvOpCopyObject:
3265 vtn_handle_composite(b, opcode, w, count);
3266 break;
3267
3268 case SpvOpEmitVertex:
3269 case SpvOpEndPrimitive:
3270 case SpvOpEmitStreamVertex:
3271 case SpvOpEndStreamPrimitive:
3272 case SpvOpControlBarrier:
3273 case SpvOpMemoryBarrier:
3274 vtn_handle_barrier(b, opcode, w, count);
3275 break;
3276
3277 default:
3278 unreachable("Unhandled opcode");
3279 }
3280
3281 return true;
3282 }
3283
3284 nir_function *
3285 spirv_to_nir(const uint32_t *words, size_t word_count,
3286 struct nir_spirv_specialization *spec, unsigned num_spec,
3287 gl_shader_stage stage, const char *entry_point_name,
3288 const struct nir_spirv_supported_extensions *ext,
3289 const nir_shader_compiler_options *options)
3290 {
3291 const uint32_t *word_end = words + word_count;
3292
3293 /* Handle the SPIR-V header (first 4 dwords) */
3294 assert(word_count > 5);
3295
3296 assert(words[0] == SpvMagicNumber);
3297 assert(words[1] >= 0x10000);
3298 /* words[2] == generator magic */
3299 unsigned value_id_bound = words[3];
3300 assert(words[4] == 0);
3301
3302 words+= 5;
3303
3304 /* Initialize the stn_builder object */
3305 struct vtn_builder *b = rzalloc(NULL, struct vtn_builder);
3306 b->value_id_bound = value_id_bound;
3307 b->values = rzalloc_array(b, struct vtn_value, value_id_bound);
3308 exec_list_make_empty(&b->functions);
3309 b->entry_point_stage = stage;
3310 b->entry_point_name = entry_point_name;
3311 b->ext = ext;
3312
3313 /* Handle all the preamble instructions */
3314 words = vtn_foreach_instruction(b, words, word_end,
3315 vtn_handle_preamble_instruction);
3316
3317 if (b->entry_point == NULL) {
3318 assert(!"Entry point not found");
3319 ralloc_free(b);
3320 return NULL;
3321 }
3322
3323 b->shader = nir_shader_create(NULL, stage, options, NULL);
3324
3325 /* Set shader info defaults */
3326 b->shader->info.gs.invocations = 1;
3327
3328 /* Parse execution modes */
3329 vtn_foreach_execution_mode(b, b->entry_point,
3330 vtn_handle_execution_mode, NULL);
3331
3332 b->specializations = spec;
3333 b->num_specializations = num_spec;
3334
3335 /* Handle all variable, type, and constant instructions */
3336 words = vtn_foreach_instruction(b, words, word_end,
3337 vtn_handle_variable_or_type_instruction);
3338
3339 vtn_build_cfg(b, words, word_end);
3340
3341 foreach_list_typed(struct vtn_function, func, node, &b->functions) {
3342 b->impl = func->impl;
3343 b->const_table = _mesa_hash_table_create(b, _mesa_hash_pointer,
3344 _mesa_key_pointer_equal);
3345
3346 vtn_function_emit(b, func, vtn_handle_body_instruction);
3347 }
3348
3349 assert(b->entry_point->value_type == vtn_value_type_function);
3350 nir_function *entry_point = b->entry_point->func->impl->function;
3351 assert(entry_point);
3352
3353 ralloc_free(b);
3354
3355 return entry_point;
3356 }