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