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