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