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