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