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