99910d2e448f396bc2f28d23e07a65fae1ebf62c
[mesa.git] / src / compiler / spirv / spirv_to_nir.c
1 /*
2 * Copyright © 2015 Intel Corporation
3 *
4 * Permission is hereby granted, free of charge, to any person obtaining a
5 * copy of this software and associated documentation files (the "Software"),
6 * to deal in the Software without restriction, including without limitation
7 * the rights to use, copy, modify, merge, publish, distribute, sublicense,
8 * and/or sell copies of the Software, and to permit persons to whom the
9 * Software is furnished to do so, subject to the following conditions:
10 *
11 * The above copyright notice and this permission notice (including the next
12 * paragraph) shall be included in all copies or substantial portions of the
13 * Software.
14 *
15 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
16 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
17 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
18 * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
19 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
20 * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
21 * IN THE SOFTWARE.
22 *
23 * Authors:
24 * Jason Ekstrand (jason@jlekstrand.net)
25 *
26 */
27
28 #include "vtn_private.h"
29 #include "nir/nir_vla.h"
30 #include "nir/nir_control_flow.h"
31 #include "nir/nir_constant_expressions.h"
32 #include "nir/nir_deref.h"
33 #include "spirv_info.h"
34
35 #include "util/format/u_format.h"
36 #include "util/u_math.h"
37
38 #include <stdio.h>
39
40 void
41 vtn_log(struct vtn_builder *b, enum nir_spirv_debug_level level,
42 size_t spirv_offset, const char *message)
43 {
44 if (b->options->debug.func) {
45 b->options->debug.func(b->options->debug.private_data,
46 level, spirv_offset, message);
47 }
48
49 #ifndef NDEBUG
50 if (level >= NIR_SPIRV_DEBUG_LEVEL_WARNING)
51 fprintf(stderr, "%s\n", message);
52 #endif
53 }
54
55 void
56 vtn_logf(struct vtn_builder *b, enum nir_spirv_debug_level level,
57 size_t spirv_offset, const char *fmt, ...)
58 {
59 va_list args;
60 char *msg;
61
62 va_start(args, fmt);
63 msg = ralloc_vasprintf(NULL, fmt, args);
64 va_end(args);
65
66 vtn_log(b, level, spirv_offset, msg);
67
68 ralloc_free(msg);
69 }
70
71 static void
72 vtn_log_err(struct vtn_builder *b,
73 enum nir_spirv_debug_level level, const char *prefix,
74 const char *file, unsigned line,
75 const char *fmt, va_list args)
76 {
77 char *msg;
78
79 msg = ralloc_strdup(NULL, prefix);
80
81 #ifndef NDEBUG
82 ralloc_asprintf_append(&msg, " In file %s:%u\n", file, line);
83 #endif
84
85 ralloc_asprintf_append(&msg, " ");
86
87 ralloc_vasprintf_append(&msg, fmt, args);
88
89 ralloc_asprintf_append(&msg, "\n %zu bytes into the SPIR-V binary",
90 b->spirv_offset);
91
92 if (b->file) {
93 ralloc_asprintf_append(&msg,
94 "\n in SPIR-V source file %s, line %d, col %d",
95 b->file, b->line, b->col);
96 }
97
98 vtn_log(b, level, b->spirv_offset, msg);
99
100 ralloc_free(msg);
101 }
102
103 static void
104 vtn_dump_shader(struct vtn_builder *b, const char *path, const char *prefix)
105 {
106 static int idx = 0;
107
108 char filename[1024];
109 int len = snprintf(filename, sizeof(filename), "%s/%s-%d.spirv",
110 path, prefix, idx++);
111 if (len < 0 || len >= sizeof(filename))
112 return;
113
114 FILE *f = fopen(filename, "w");
115 if (f == NULL)
116 return;
117
118 fwrite(b->spirv, sizeof(*b->spirv), b->spirv_word_count, f);
119 fclose(f);
120
121 vtn_info("SPIR-V shader dumped to %s", filename);
122 }
123
124 void
125 _vtn_warn(struct vtn_builder *b, const char *file, unsigned line,
126 const char *fmt, ...)
127 {
128 va_list args;
129
130 va_start(args, fmt);
131 vtn_log_err(b, NIR_SPIRV_DEBUG_LEVEL_WARNING, "SPIR-V WARNING:\n",
132 file, line, fmt, args);
133 va_end(args);
134 }
135
136 void
137 _vtn_err(struct vtn_builder *b, const char *file, unsigned line,
138 const char *fmt, ...)
139 {
140 va_list args;
141
142 va_start(args, fmt);
143 vtn_log_err(b, NIR_SPIRV_DEBUG_LEVEL_ERROR, "SPIR-V ERROR:\n",
144 file, line, fmt, args);
145 va_end(args);
146 }
147
148 void
149 _vtn_fail(struct vtn_builder *b, const char *file, unsigned line,
150 const char *fmt, ...)
151 {
152 va_list args;
153
154 va_start(args, fmt);
155 vtn_log_err(b, NIR_SPIRV_DEBUG_LEVEL_ERROR, "SPIR-V parsing FAILED:\n",
156 file, line, fmt, args);
157 va_end(args);
158
159 const char *dump_path = getenv("MESA_SPIRV_FAIL_DUMP_PATH");
160 if (dump_path)
161 vtn_dump_shader(b, dump_path, "fail");
162
163 longjmp(b->fail_jump, 1);
164 }
165
166 static struct vtn_ssa_value *
167 vtn_undef_ssa_value(struct vtn_builder *b, const struct glsl_type *type)
168 {
169 struct vtn_ssa_value *val = rzalloc(b, struct vtn_ssa_value);
170 val->type = glsl_get_bare_type(type);
171
172 if (glsl_type_is_vector_or_scalar(type)) {
173 unsigned num_components = glsl_get_vector_elements(val->type);
174 unsigned bit_size = glsl_get_bit_size(val->type);
175 val->def = nir_ssa_undef(&b->nb, num_components, bit_size);
176 } else {
177 unsigned elems = glsl_get_length(val->type);
178 val->elems = ralloc_array(b, struct vtn_ssa_value *, elems);
179 if (glsl_type_is_array_or_matrix(type)) {
180 const struct glsl_type *elem_type = glsl_get_array_element(type);
181 for (unsigned i = 0; i < elems; i++)
182 val->elems[i] = vtn_undef_ssa_value(b, elem_type);
183 } else {
184 vtn_assert(glsl_type_is_struct_or_ifc(type));
185 for (unsigned i = 0; i < elems; i++) {
186 const struct glsl_type *elem_type = glsl_get_struct_field(type, i);
187 val->elems[i] = vtn_undef_ssa_value(b, elem_type);
188 }
189 }
190 }
191
192 return val;
193 }
194
195 static struct vtn_ssa_value *
196 vtn_const_ssa_value(struct vtn_builder *b, nir_constant *constant,
197 const struct glsl_type *type)
198 {
199 struct hash_entry *entry = _mesa_hash_table_search(b->const_table, constant);
200
201 if (entry)
202 return entry->data;
203
204 struct vtn_ssa_value *val = rzalloc(b, struct vtn_ssa_value);
205 val->type = glsl_get_bare_type(type);
206
207 if (glsl_type_is_vector_or_scalar(type)) {
208 unsigned num_components = glsl_get_vector_elements(val->type);
209 unsigned bit_size = glsl_get_bit_size(type);
210 nir_load_const_instr *load =
211 nir_load_const_instr_create(b->shader, num_components, bit_size);
212
213 memcpy(load->value, constant->values,
214 sizeof(nir_const_value) * num_components);
215
216 nir_instr_insert_before_cf_list(&b->nb.impl->body, &load->instr);
217 val->def = &load->def;
218 } else {
219 unsigned elems = glsl_get_length(val->type);
220 val->elems = ralloc_array(b, struct vtn_ssa_value *, elems);
221 if (glsl_type_is_array_or_matrix(type)) {
222 const struct glsl_type *elem_type = glsl_get_array_element(type);
223 for (unsigned i = 0; i < elems; i++) {
224 val->elems[i] = vtn_const_ssa_value(b, constant->elements[i],
225 elem_type);
226 }
227 } else {
228 vtn_assert(glsl_type_is_struct_or_ifc(type));
229 for (unsigned i = 0; i < elems; i++) {
230 const struct glsl_type *elem_type = glsl_get_struct_field(type, i);
231 val->elems[i] = vtn_const_ssa_value(b, constant->elements[i],
232 elem_type);
233 }
234 }
235 }
236
237 return val;
238 }
239
240 struct vtn_ssa_value *
241 vtn_ssa_value(struct vtn_builder *b, uint32_t value_id)
242 {
243 struct vtn_value *val = vtn_untyped_value(b, value_id);
244 switch (val->value_type) {
245 case vtn_value_type_undef:
246 return vtn_undef_ssa_value(b, val->type->type);
247
248 case vtn_value_type_constant:
249 return vtn_const_ssa_value(b, val->constant, val->type->type);
250
251 case vtn_value_type_ssa:
252 return val->ssa;
253
254 case vtn_value_type_pointer:
255 vtn_assert(val->pointer->ptr_type && val->pointer->ptr_type->type);
256 struct vtn_ssa_value *ssa =
257 vtn_create_ssa_value(b, val->pointer->ptr_type->type);
258 ssa->def = vtn_pointer_to_ssa(b, val->pointer);
259 return ssa;
260
261 default:
262 vtn_fail("Invalid type for an SSA value");
263 }
264 }
265
266 struct vtn_value *
267 vtn_push_ssa_value(struct vtn_builder *b, uint32_t value_id,
268 struct vtn_ssa_value *ssa)
269 {
270 struct vtn_type *type = vtn_get_value_type(b, value_id);
271
272 /* See vtn_create_ssa_value */
273 vtn_fail_if(ssa->type != glsl_get_bare_type(type->type),
274 "Type mismatch for SPIR-V SSA value");
275
276 struct vtn_value *val;
277 if (type->base_type == vtn_base_type_pointer) {
278 val = vtn_push_pointer(b, value_id, vtn_pointer_from_ssa(b, ssa->def, type));
279 } else {
280 /* Don't trip the value_type_ssa check in vtn_push_value */
281 val = vtn_push_value(b, value_id, vtn_value_type_invalid);
282 val->value_type = vtn_value_type_ssa;
283 val->ssa = ssa;
284 }
285
286 return val;
287 }
288
289 nir_ssa_def *
290 vtn_get_nir_ssa(struct vtn_builder *b, uint32_t value_id)
291 {
292 struct vtn_ssa_value *ssa = vtn_ssa_value(b, value_id);
293 vtn_fail_if(!glsl_type_is_vector_or_scalar(ssa->type),
294 "Expected a vector or scalar type");
295 return ssa->def;
296 }
297
298 struct vtn_value *
299 vtn_push_nir_ssa(struct vtn_builder *b, uint32_t value_id, nir_ssa_def *def)
300 {
301 /* Types for all SPIR-V SSA values are set as part of a pre-pass so the
302 * type will be valid by the time we get here.
303 */
304 struct vtn_type *type = vtn_get_value_type(b, value_id);
305 vtn_fail_if(def->num_components != glsl_get_vector_elements(type->type) ||
306 def->bit_size != glsl_get_bit_size(type->type),
307 "Mismatch between NIR and SPIR-V type.");
308 struct vtn_ssa_value *ssa = vtn_create_ssa_value(b, type->type);
309 ssa->def = def;
310 return vtn_push_ssa_value(b, value_id, ssa);
311 }
312
313 static nir_deref_instr *
314 vtn_get_image(struct vtn_builder *b, uint32_t value_id)
315 {
316 struct vtn_type *type = vtn_get_value_type(b, value_id);
317 vtn_assert(type->base_type == vtn_base_type_image);
318 return nir_build_deref_cast(&b->nb, vtn_get_nir_ssa(b, value_id),
319 nir_var_uniform, type->glsl_image, 0);
320 }
321
322 static void
323 vtn_push_image(struct vtn_builder *b, uint32_t value_id,
324 nir_deref_instr *deref)
325 {
326 struct vtn_type *type = vtn_get_value_type(b, value_id);
327 vtn_assert(type->base_type == vtn_base_type_image);
328 vtn_push_nir_ssa(b, value_id, &deref->dest.ssa);
329 }
330
331 static nir_deref_instr *
332 vtn_get_sampler(struct vtn_builder *b, uint32_t value_id)
333 {
334 struct vtn_type *type = vtn_get_value_type(b, value_id);
335 vtn_assert(type->base_type == vtn_base_type_sampler);
336 return nir_build_deref_cast(&b->nb, vtn_get_nir_ssa(b, value_id),
337 nir_var_uniform, glsl_bare_sampler_type(), 0);
338 }
339
340 nir_ssa_def *
341 vtn_sampled_image_to_nir_ssa(struct vtn_builder *b,
342 struct vtn_sampled_image si)
343 {
344 return nir_vec2(&b->nb, &si.image->dest.ssa, &si.sampler->dest.ssa);
345 }
346
347 static void
348 vtn_push_sampled_image(struct vtn_builder *b, uint32_t value_id,
349 struct vtn_sampled_image si)
350 {
351 struct vtn_type *type = vtn_get_value_type(b, value_id);
352 vtn_assert(type->base_type == vtn_base_type_sampled_image);
353 vtn_push_nir_ssa(b, value_id, vtn_sampled_image_to_nir_ssa(b, si));
354 }
355
356 static struct vtn_sampled_image
357 vtn_get_sampled_image(struct vtn_builder *b, uint32_t value_id)
358 {
359 struct vtn_type *type = vtn_get_value_type(b, value_id);
360 vtn_assert(type->base_type == vtn_base_type_sampled_image);
361 nir_ssa_def *si_vec2 = vtn_get_nir_ssa(b, value_id);
362
363 struct vtn_sampled_image si = { NULL, };
364 si.image = nir_build_deref_cast(&b->nb, nir_channel(&b->nb, si_vec2, 0),
365 nir_var_uniform,
366 type->image->glsl_image, 0);
367 si.sampler = nir_build_deref_cast(&b->nb, nir_channel(&b->nb, si_vec2, 1),
368 nir_var_uniform,
369 glsl_bare_sampler_type(), 0);
370 return si;
371 }
372
373 static char *
374 vtn_string_literal(struct vtn_builder *b, const uint32_t *words,
375 unsigned word_count, unsigned *words_used)
376 {
377 char *dup = ralloc_strndup(b, (char *)words, word_count * sizeof(*words));
378 if (words_used) {
379 /* Ammount of space taken by the string (including the null) */
380 unsigned len = strlen(dup) + 1;
381 *words_used = DIV_ROUND_UP(len, sizeof(*words));
382 }
383 return dup;
384 }
385
386 const uint32_t *
387 vtn_foreach_instruction(struct vtn_builder *b, const uint32_t *start,
388 const uint32_t *end, vtn_instruction_handler handler)
389 {
390 b->file = NULL;
391 b->line = -1;
392 b->col = -1;
393
394 const uint32_t *w = start;
395 while (w < end) {
396 SpvOp opcode = w[0] & SpvOpCodeMask;
397 unsigned count = w[0] >> SpvWordCountShift;
398 vtn_assert(count >= 1 && w + count <= end);
399
400 b->spirv_offset = (uint8_t *)w - (uint8_t *)b->spirv;
401
402 switch (opcode) {
403 case SpvOpNop:
404 break; /* Do nothing */
405
406 case SpvOpLine:
407 b->file = vtn_value(b, w[1], vtn_value_type_string)->str;
408 b->line = w[2];
409 b->col = w[3];
410 break;
411
412 case SpvOpNoLine:
413 b->file = NULL;
414 b->line = -1;
415 b->col = -1;
416 break;
417
418 default:
419 if (!handler(b, opcode, w, count))
420 return w;
421 break;
422 }
423
424 w += count;
425 }
426
427 b->spirv_offset = 0;
428 b->file = NULL;
429 b->line = -1;
430 b->col = -1;
431
432 assert(w == end);
433 return w;
434 }
435
436 static bool
437 vtn_handle_non_semantic_instruction(struct vtn_builder *b, SpvOp ext_opcode,
438 const uint32_t *w, unsigned count)
439 {
440 /* Do nothing. */
441 return true;
442 }
443
444 static void
445 vtn_handle_extension(struct vtn_builder *b, SpvOp opcode,
446 const uint32_t *w, unsigned count)
447 {
448 const char *ext = (const char *)&w[2];
449 switch (opcode) {
450 case SpvOpExtInstImport: {
451 struct vtn_value *val = vtn_push_value(b, w[1], vtn_value_type_extension);
452 if (strcmp(ext, "GLSL.std.450") == 0) {
453 val->ext_handler = vtn_handle_glsl450_instruction;
454 } else if ((strcmp(ext, "SPV_AMD_gcn_shader") == 0)
455 && (b->options && b->options->caps.amd_gcn_shader)) {
456 val->ext_handler = vtn_handle_amd_gcn_shader_instruction;
457 } else if ((strcmp(ext, "SPV_AMD_shader_ballot") == 0)
458 && (b->options && b->options->caps.amd_shader_ballot)) {
459 val->ext_handler = vtn_handle_amd_shader_ballot_instruction;
460 } else if ((strcmp(ext, "SPV_AMD_shader_trinary_minmax") == 0)
461 && (b->options && b->options->caps.amd_trinary_minmax)) {
462 val->ext_handler = vtn_handle_amd_shader_trinary_minmax_instruction;
463 } else if ((strcmp(ext, "SPV_AMD_shader_explicit_vertex_parameter") == 0)
464 && (b->options && b->options->caps.amd_shader_explicit_vertex_parameter)) {
465 val->ext_handler = vtn_handle_amd_shader_explicit_vertex_parameter_instruction;
466 } else if (strcmp(ext, "OpenCL.std") == 0) {
467 val->ext_handler = vtn_handle_opencl_instruction;
468 } else if (strstr(ext, "NonSemantic.") == ext) {
469 val->ext_handler = vtn_handle_non_semantic_instruction;
470 } else {
471 vtn_fail("Unsupported extension: %s", ext);
472 }
473 break;
474 }
475
476 case SpvOpExtInst: {
477 struct vtn_value *val = vtn_value(b, w[3], vtn_value_type_extension);
478 bool handled = val->ext_handler(b, w[4], w, count);
479 vtn_assert(handled);
480 break;
481 }
482
483 default:
484 vtn_fail_with_opcode("Unhandled opcode", opcode);
485 }
486 }
487
488 static void
489 _foreach_decoration_helper(struct vtn_builder *b,
490 struct vtn_value *base_value,
491 int parent_member,
492 struct vtn_value *value,
493 vtn_decoration_foreach_cb cb, void *data)
494 {
495 for (struct vtn_decoration *dec = value->decoration; dec; dec = dec->next) {
496 int member;
497 if (dec->scope == VTN_DEC_DECORATION) {
498 member = parent_member;
499 } else if (dec->scope >= VTN_DEC_STRUCT_MEMBER0) {
500 vtn_fail_if(value->value_type != vtn_value_type_type ||
501 value->type->base_type != vtn_base_type_struct,
502 "OpMemberDecorate and OpGroupMemberDecorate are only "
503 "allowed on OpTypeStruct");
504 /* This means we haven't recursed yet */
505 assert(value == base_value);
506
507 member = dec->scope - VTN_DEC_STRUCT_MEMBER0;
508
509 vtn_fail_if(member >= base_value->type->length,
510 "OpMemberDecorate specifies member %d but the "
511 "OpTypeStruct has only %u members",
512 member, base_value->type->length);
513 } else {
514 /* Not a decoration */
515 assert(dec->scope == VTN_DEC_EXECUTION_MODE);
516 continue;
517 }
518
519 if (dec->group) {
520 assert(dec->group->value_type == vtn_value_type_decoration_group);
521 _foreach_decoration_helper(b, base_value, member, dec->group,
522 cb, data);
523 } else {
524 cb(b, base_value, member, dec, data);
525 }
526 }
527 }
528
529 /** Iterates (recursively if needed) over all of the decorations on a value
530 *
531 * This function iterates over all of the decorations applied to a given
532 * value. If it encounters a decoration group, it recurses into the group
533 * and iterates over all of those decorations as well.
534 */
535 void
536 vtn_foreach_decoration(struct vtn_builder *b, struct vtn_value *value,
537 vtn_decoration_foreach_cb cb, void *data)
538 {
539 _foreach_decoration_helper(b, value, -1, value, cb, data);
540 }
541
542 void
543 vtn_foreach_execution_mode(struct vtn_builder *b, struct vtn_value *value,
544 vtn_execution_mode_foreach_cb cb, void *data)
545 {
546 for (struct vtn_decoration *dec = value->decoration; dec; dec = dec->next) {
547 if (dec->scope != VTN_DEC_EXECUTION_MODE)
548 continue;
549
550 assert(dec->group == NULL);
551 cb(b, value, dec, data);
552 }
553 }
554
555 void
556 vtn_handle_decoration(struct vtn_builder *b, SpvOp opcode,
557 const uint32_t *w, unsigned count)
558 {
559 const uint32_t *w_end = w + count;
560 const uint32_t target = w[1];
561 w += 2;
562
563 switch (opcode) {
564 case SpvOpDecorationGroup:
565 vtn_push_value(b, target, vtn_value_type_decoration_group);
566 break;
567
568 case SpvOpDecorate:
569 case SpvOpDecorateId:
570 case SpvOpMemberDecorate:
571 case SpvOpDecorateString:
572 case SpvOpMemberDecorateString:
573 case SpvOpExecutionMode:
574 case SpvOpExecutionModeId: {
575 struct vtn_value *val = vtn_untyped_value(b, target);
576
577 struct vtn_decoration *dec = rzalloc(b, struct vtn_decoration);
578 switch (opcode) {
579 case SpvOpDecorate:
580 case SpvOpDecorateId:
581 case SpvOpDecorateString:
582 dec->scope = VTN_DEC_DECORATION;
583 break;
584 case SpvOpMemberDecorate:
585 case SpvOpMemberDecorateString:
586 dec->scope = VTN_DEC_STRUCT_MEMBER0 + *(w++);
587 vtn_fail_if(dec->scope < VTN_DEC_STRUCT_MEMBER0, /* overflow */
588 "Member argument of OpMemberDecorate too large");
589 break;
590 case SpvOpExecutionMode:
591 case SpvOpExecutionModeId:
592 dec->scope = VTN_DEC_EXECUTION_MODE;
593 break;
594 default:
595 unreachable("Invalid decoration opcode");
596 }
597 dec->decoration = *(w++);
598 dec->operands = w;
599
600 /* Link into the list */
601 dec->next = val->decoration;
602 val->decoration = dec;
603 break;
604 }
605
606 case SpvOpGroupMemberDecorate:
607 case SpvOpGroupDecorate: {
608 struct vtn_value *group =
609 vtn_value(b, target, vtn_value_type_decoration_group);
610
611 for (; w < w_end; w++) {
612 struct vtn_value *val = vtn_untyped_value(b, *w);
613 struct vtn_decoration *dec = rzalloc(b, struct vtn_decoration);
614
615 dec->group = group;
616 if (opcode == SpvOpGroupDecorate) {
617 dec->scope = VTN_DEC_DECORATION;
618 } else {
619 dec->scope = VTN_DEC_STRUCT_MEMBER0 + *(++w);
620 vtn_fail_if(dec->scope < 0, /* Check for overflow */
621 "Member argument of OpGroupMemberDecorate too large");
622 }
623
624 /* Link into the list */
625 dec->next = val->decoration;
626 val->decoration = dec;
627 }
628 break;
629 }
630
631 default:
632 unreachable("Unhandled opcode");
633 }
634 }
635
636 struct member_decoration_ctx {
637 unsigned num_fields;
638 struct glsl_struct_field *fields;
639 struct vtn_type *type;
640 };
641
642 /**
643 * Returns true if the given type contains a struct decorated Block or
644 * BufferBlock
645 */
646 bool
647 vtn_type_contains_block(struct vtn_builder *b, struct vtn_type *type)
648 {
649 switch (type->base_type) {
650 case vtn_base_type_array:
651 return vtn_type_contains_block(b, type->array_element);
652 case vtn_base_type_struct:
653 if (type->block || type->buffer_block)
654 return true;
655 for (unsigned i = 0; i < type->length; i++) {
656 if (vtn_type_contains_block(b, type->members[i]))
657 return true;
658 }
659 return false;
660 default:
661 return false;
662 }
663 }
664
665 /** Returns true if two types are "compatible", i.e. you can do an OpLoad,
666 * OpStore, or OpCopyMemory between them without breaking anything.
667 * Technically, the SPIR-V rules require the exact same type ID but this lets
668 * us internally be a bit looser.
669 */
670 bool
671 vtn_types_compatible(struct vtn_builder *b,
672 struct vtn_type *t1, struct vtn_type *t2)
673 {
674 if (t1->id == t2->id)
675 return true;
676
677 if (t1->base_type != t2->base_type)
678 return false;
679
680 switch (t1->base_type) {
681 case vtn_base_type_void:
682 case vtn_base_type_scalar:
683 case vtn_base_type_vector:
684 case vtn_base_type_matrix:
685 case vtn_base_type_image:
686 case vtn_base_type_sampler:
687 case vtn_base_type_sampled_image:
688 return t1->type == t2->type;
689
690 case vtn_base_type_array:
691 return t1->length == t2->length &&
692 vtn_types_compatible(b, t1->array_element, t2->array_element);
693
694 case vtn_base_type_pointer:
695 return vtn_types_compatible(b, t1->deref, t2->deref);
696
697 case vtn_base_type_struct:
698 if (t1->length != t2->length)
699 return false;
700
701 for (unsigned i = 0; i < t1->length; i++) {
702 if (!vtn_types_compatible(b, t1->members[i], t2->members[i]))
703 return false;
704 }
705 return true;
706
707 case vtn_base_type_function:
708 /* This case shouldn't get hit since you can't copy around function
709 * types. Just require them to be identical.
710 */
711 return false;
712 }
713
714 vtn_fail("Invalid base type");
715 }
716
717 struct vtn_type *
718 vtn_type_without_array(struct vtn_type *type)
719 {
720 while (type->base_type == vtn_base_type_array)
721 type = type->array_element;
722 return type;
723 }
724
725 /* does a shallow copy of a vtn_type */
726
727 static struct vtn_type *
728 vtn_type_copy(struct vtn_builder *b, struct vtn_type *src)
729 {
730 struct vtn_type *dest = ralloc(b, struct vtn_type);
731 *dest = *src;
732
733 switch (src->base_type) {
734 case vtn_base_type_void:
735 case vtn_base_type_scalar:
736 case vtn_base_type_vector:
737 case vtn_base_type_matrix:
738 case vtn_base_type_array:
739 case vtn_base_type_pointer:
740 case vtn_base_type_image:
741 case vtn_base_type_sampler:
742 case vtn_base_type_sampled_image:
743 /* Nothing more to do */
744 break;
745
746 case vtn_base_type_struct:
747 dest->members = ralloc_array(b, struct vtn_type *, src->length);
748 memcpy(dest->members, src->members,
749 src->length * sizeof(src->members[0]));
750
751 dest->offsets = ralloc_array(b, unsigned, src->length);
752 memcpy(dest->offsets, src->offsets,
753 src->length * sizeof(src->offsets[0]));
754 break;
755
756 case vtn_base_type_function:
757 dest->params = ralloc_array(b, struct vtn_type *, src->length);
758 memcpy(dest->params, src->params, src->length * sizeof(src->params[0]));
759 break;
760 }
761
762 return dest;
763 }
764
765 static const struct glsl_type *
766 wrap_type_in_array(const struct glsl_type *type,
767 const struct glsl_type *array_type)
768 {
769 if (!glsl_type_is_array(array_type))
770 return type;
771
772 const struct glsl_type *elem_type =
773 wrap_type_in_array(type, glsl_get_array_element(array_type));
774 return glsl_array_type(elem_type, glsl_get_length(array_type),
775 glsl_get_explicit_stride(array_type));
776 }
777
778 static bool
779 vtn_type_needs_explicit_layout(struct vtn_builder *b, enum vtn_variable_mode mode)
780 {
781 /* For OpenCL we never want to strip the info from the types, and it makes
782 * type comparisons easier in later stages.
783 */
784 if (b->options->environment == NIR_SPIRV_OPENCL)
785 return true;
786
787 switch (mode) {
788 case vtn_variable_mode_input:
789 case vtn_variable_mode_output:
790 /* Layout decorations kept because we need offsets for XFB arrays of
791 * blocks.
792 */
793 return b->shader->info.has_transform_feedback_varyings;
794
795 case vtn_variable_mode_ssbo:
796 case vtn_variable_mode_phys_ssbo:
797 case vtn_variable_mode_ubo:
798 return true;
799
800 default:
801 return false;
802 }
803 }
804
805 const struct glsl_type *
806 vtn_type_get_nir_type(struct vtn_builder *b, struct vtn_type *type,
807 enum vtn_variable_mode mode)
808 {
809 if (mode == vtn_variable_mode_atomic_counter) {
810 vtn_fail_if(glsl_without_array(type->type) != glsl_uint_type(),
811 "Variables in the AtomicCounter storage class should be "
812 "(possibly arrays of arrays of) uint.");
813 return wrap_type_in_array(glsl_atomic_uint_type(), type->type);
814 }
815
816 if (mode == vtn_variable_mode_uniform) {
817 switch (type->base_type) {
818 case vtn_base_type_array: {
819 const struct glsl_type *elem_type =
820 vtn_type_get_nir_type(b, type->array_element, mode);
821
822 return glsl_array_type(elem_type, type->length,
823 glsl_get_explicit_stride(type->type));
824 }
825
826 case vtn_base_type_struct: {
827 bool need_new_struct = false;
828 const uint32_t num_fields = type->length;
829 NIR_VLA(struct glsl_struct_field, fields, num_fields);
830 for (unsigned i = 0; i < num_fields; i++) {
831 fields[i] = *glsl_get_struct_field_data(type->type, i);
832 const struct glsl_type *field_nir_type =
833 vtn_type_get_nir_type(b, type->members[i], mode);
834 if (fields[i].type != field_nir_type) {
835 fields[i].type = field_nir_type;
836 need_new_struct = true;
837 }
838 }
839 if (need_new_struct) {
840 if (glsl_type_is_interface(type->type)) {
841 return glsl_interface_type(fields, num_fields,
842 /* packing */ 0, false,
843 glsl_get_type_name(type->type));
844 } else {
845 return glsl_struct_type(fields, num_fields,
846 glsl_get_type_name(type->type),
847 glsl_struct_type_is_packed(type->type));
848 }
849 } else {
850 /* No changes, just pass it on */
851 return type->type;
852 }
853 }
854
855 case vtn_base_type_image:
856 return type->glsl_image;
857
858 case vtn_base_type_sampler:
859 return glsl_bare_sampler_type();
860
861 case vtn_base_type_sampled_image:
862 return type->image->glsl_image;
863
864 default:
865 return type->type;
866 }
867 }
868
869 /* Layout decorations are allowed but ignored in certain conditions,
870 * to allow SPIR-V generators perform type deduplication. Discard
871 * unnecessary ones when passing to NIR.
872 */
873 if (!vtn_type_needs_explicit_layout(b, mode))
874 return glsl_get_bare_type(type->type);
875
876 return type->type;
877 }
878
879 static struct vtn_type *
880 mutable_matrix_member(struct vtn_builder *b, struct vtn_type *type, int member)
881 {
882 type->members[member] = vtn_type_copy(b, type->members[member]);
883 type = type->members[member];
884
885 /* We may have an array of matrices.... Oh, joy! */
886 while (glsl_type_is_array(type->type)) {
887 type->array_element = vtn_type_copy(b, type->array_element);
888 type = type->array_element;
889 }
890
891 vtn_assert(glsl_type_is_matrix(type->type));
892
893 return type;
894 }
895
896 static void
897 vtn_handle_access_qualifier(struct vtn_builder *b, struct vtn_type *type,
898 int member, enum gl_access_qualifier access)
899 {
900 type->members[member] = vtn_type_copy(b, type->members[member]);
901 type = type->members[member];
902
903 type->access |= access;
904 }
905
906 static void
907 array_stride_decoration_cb(struct vtn_builder *b,
908 struct vtn_value *val, int member,
909 const struct vtn_decoration *dec, void *void_ctx)
910 {
911 struct vtn_type *type = val->type;
912
913 if (dec->decoration == SpvDecorationArrayStride) {
914 if (vtn_type_contains_block(b, type)) {
915 vtn_warn("The ArrayStride decoration cannot be applied to an array "
916 "type which contains a structure type decorated Block "
917 "or BufferBlock");
918 /* Ignore the decoration */
919 } else {
920 vtn_fail_if(dec->operands[0] == 0, "ArrayStride must be non-zero");
921 type->stride = dec->operands[0];
922 }
923 }
924 }
925
926 static void
927 struct_member_decoration_cb(struct vtn_builder *b,
928 UNUSED struct vtn_value *val, int member,
929 const struct vtn_decoration *dec, void *void_ctx)
930 {
931 struct member_decoration_ctx *ctx = void_ctx;
932
933 if (member < 0)
934 return;
935
936 assert(member < ctx->num_fields);
937
938 switch (dec->decoration) {
939 case SpvDecorationRelaxedPrecision:
940 case SpvDecorationUniform:
941 case SpvDecorationUniformId:
942 break; /* FIXME: Do nothing with this for now. */
943 case SpvDecorationNonWritable:
944 vtn_handle_access_qualifier(b, ctx->type, member, ACCESS_NON_WRITEABLE);
945 break;
946 case SpvDecorationNonReadable:
947 vtn_handle_access_qualifier(b, ctx->type, member, ACCESS_NON_READABLE);
948 break;
949 case SpvDecorationVolatile:
950 vtn_handle_access_qualifier(b, ctx->type, member, ACCESS_VOLATILE);
951 break;
952 case SpvDecorationCoherent:
953 vtn_handle_access_qualifier(b, ctx->type, member, ACCESS_COHERENT);
954 break;
955 case SpvDecorationNoPerspective:
956 ctx->fields[member].interpolation = INTERP_MODE_NOPERSPECTIVE;
957 break;
958 case SpvDecorationFlat:
959 ctx->fields[member].interpolation = INTERP_MODE_FLAT;
960 break;
961 case SpvDecorationExplicitInterpAMD:
962 ctx->fields[member].interpolation = INTERP_MODE_EXPLICIT;
963 break;
964 case SpvDecorationCentroid:
965 ctx->fields[member].centroid = true;
966 break;
967 case SpvDecorationSample:
968 ctx->fields[member].sample = true;
969 break;
970 case SpvDecorationStream:
971 /* This is handled later by var_decoration_cb in vtn_variables.c */
972 break;
973 case SpvDecorationLocation:
974 ctx->fields[member].location = dec->operands[0];
975 break;
976 case SpvDecorationComponent:
977 break; /* FIXME: What should we do with these? */
978 case SpvDecorationBuiltIn:
979 ctx->type->members[member] = vtn_type_copy(b, ctx->type->members[member]);
980 ctx->type->members[member]->is_builtin = true;
981 ctx->type->members[member]->builtin = dec->operands[0];
982 ctx->type->builtin_block = true;
983 break;
984 case SpvDecorationOffset:
985 ctx->type->offsets[member] = dec->operands[0];
986 ctx->fields[member].offset = dec->operands[0];
987 break;
988 case SpvDecorationMatrixStride:
989 /* Handled as a second pass */
990 break;
991 case SpvDecorationColMajor:
992 break; /* Nothing to do here. Column-major is the default. */
993 case SpvDecorationRowMajor:
994 mutable_matrix_member(b, ctx->type, member)->row_major = true;
995 break;
996
997 case SpvDecorationPatch:
998 break;
999
1000 case SpvDecorationSpecId:
1001 case SpvDecorationBlock:
1002 case SpvDecorationBufferBlock:
1003 case SpvDecorationArrayStride:
1004 case SpvDecorationGLSLShared:
1005 case SpvDecorationGLSLPacked:
1006 case SpvDecorationInvariant:
1007 case SpvDecorationRestrict:
1008 case SpvDecorationAliased:
1009 case SpvDecorationConstant:
1010 case SpvDecorationIndex:
1011 case SpvDecorationBinding:
1012 case SpvDecorationDescriptorSet:
1013 case SpvDecorationLinkageAttributes:
1014 case SpvDecorationNoContraction:
1015 case SpvDecorationInputAttachmentIndex:
1016 vtn_warn("Decoration not allowed on struct members: %s",
1017 spirv_decoration_to_string(dec->decoration));
1018 break;
1019
1020 case SpvDecorationXfbBuffer:
1021 case SpvDecorationXfbStride:
1022 /* This is handled later by var_decoration_cb in vtn_variables.c */
1023 break;
1024
1025 case SpvDecorationCPacked:
1026 if (b->shader->info.stage != MESA_SHADER_KERNEL)
1027 vtn_warn("Decoration only allowed for CL-style kernels: %s",
1028 spirv_decoration_to_string(dec->decoration));
1029 else
1030 ctx->type->packed = true;
1031 break;
1032
1033 case SpvDecorationSaturatedConversion:
1034 case SpvDecorationFuncParamAttr:
1035 case SpvDecorationFPRoundingMode:
1036 case SpvDecorationFPFastMathMode:
1037 case SpvDecorationAlignment:
1038 if (b->shader->info.stage != MESA_SHADER_KERNEL) {
1039 vtn_warn("Decoration only allowed for CL-style kernels: %s",
1040 spirv_decoration_to_string(dec->decoration));
1041 }
1042 break;
1043
1044 case SpvDecorationUserSemantic:
1045 case SpvDecorationUserTypeGOOGLE:
1046 /* User semantic decorations can safely be ignored by the driver. */
1047 break;
1048
1049 default:
1050 vtn_fail_with_decoration("Unhandled decoration", dec->decoration);
1051 }
1052 }
1053
1054 /** Chases the array type all the way down to the tail and rewrites the
1055 * glsl_types to be based off the tail's glsl_type.
1056 */
1057 static void
1058 vtn_array_type_rewrite_glsl_type(struct vtn_type *type)
1059 {
1060 if (type->base_type != vtn_base_type_array)
1061 return;
1062
1063 vtn_array_type_rewrite_glsl_type(type->array_element);
1064
1065 type->type = glsl_array_type(type->array_element->type,
1066 type->length, type->stride);
1067 }
1068
1069 /* Matrix strides are handled as a separate pass because we need to know
1070 * whether the matrix is row-major or not first.
1071 */
1072 static void
1073 struct_member_matrix_stride_cb(struct vtn_builder *b,
1074 UNUSED struct vtn_value *val, int member,
1075 const struct vtn_decoration *dec,
1076 void *void_ctx)
1077 {
1078 if (dec->decoration != SpvDecorationMatrixStride)
1079 return;
1080
1081 vtn_fail_if(member < 0,
1082 "The MatrixStride decoration is only allowed on members "
1083 "of OpTypeStruct");
1084 vtn_fail_if(dec->operands[0] == 0, "MatrixStride must be non-zero");
1085
1086 struct member_decoration_ctx *ctx = void_ctx;
1087
1088 struct vtn_type *mat_type = mutable_matrix_member(b, ctx->type, member);
1089 if (mat_type->row_major) {
1090 mat_type->array_element = vtn_type_copy(b, mat_type->array_element);
1091 mat_type->stride = mat_type->array_element->stride;
1092 mat_type->array_element->stride = dec->operands[0];
1093
1094 mat_type->type = glsl_explicit_matrix_type(mat_type->type,
1095 dec->operands[0], true);
1096 mat_type->array_element->type = glsl_get_column_type(mat_type->type);
1097 } else {
1098 vtn_assert(mat_type->array_element->stride > 0);
1099 mat_type->stride = dec->operands[0];
1100
1101 mat_type->type = glsl_explicit_matrix_type(mat_type->type,
1102 dec->operands[0], false);
1103 }
1104
1105 /* Now that we've replaced the glsl_type with a properly strided matrix
1106 * type, rewrite the member type so that it's an array of the proper kind
1107 * of glsl_type.
1108 */
1109 vtn_array_type_rewrite_glsl_type(ctx->type->members[member]);
1110 ctx->fields[member].type = ctx->type->members[member]->type;
1111 }
1112
1113 static void
1114 struct_block_decoration_cb(struct vtn_builder *b,
1115 struct vtn_value *val, int member,
1116 const struct vtn_decoration *dec, void *ctx)
1117 {
1118 if (member != -1)
1119 return;
1120
1121 struct vtn_type *type = val->type;
1122 if (dec->decoration == SpvDecorationBlock)
1123 type->block = true;
1124 else if (dec->decoration == SpvDecorationBufferBlock)
1125 type->buffer_block = true;
1126 }
1127
1128 static void
1129 type_decoration_cb(struct vtn_builder *b,
1130 struct vtn_value *val, int member,
1131 const struct vtn_decoration *dec, UNUSED void *ctx)
1132 {
1133 struct vtn_type *type = val->type;
1134
1135 if (member != -1) {
1136 /* This should have been handled by OpTypeStruct */
1137 assert(val->type->base_type == vtn_base_type_struct);
1138 assert(member >= 0 && member < val->type->length);
1139 return;
1140 }
1141
1142 switch (dec->decoration) {
1143 case SpvDecorationArrayStride:
1144 vtn_assert(type->base_type == vtn_base_type_array ||
1145 type->base_type == vtn_base_type_pointer);
1146 break;
1147 case SpvDecorationBlock:
1148 vtn_assert(type->base_type == vtn_base_type_struct);
1149 vtn_assert(type->block);
1150 break;
1151 case SpvDecorationBufferBlock:
1152 vtn_assert(type->base_type == vtn_base_type_struct);
1153 vtn_assert(type->buffer_block);
1154 break;
1155 case SpvDecorationGLSLShared:
1156 case SpvDecorationGLSLPacked:
1157 /* Ignore these, since we get explicit offsets anyways */
1158 break;
1159
1160 case SpvDecorationRowMajor:
1161 case SpvDecorationColMajor:
1162 case SpvDecorationMatrixStride:
1163 case SpvDecorationBuiltIn:
1164 case SpvDecorationNoPerspective:
1165 case SpvDecorationFlat:
1166 case SpvDecorationPatch:
1167 case SpvDecorationCentroid:
1168 case SpvDecorationSample:
1169 case SpvDecorationExplicitInterpAMD:
1170 case SpvDecorationVolatile:
1171 case SpvDecorationCoherent:
1172 case SpvDecorationNonWritable:
1173 case SpvDecorationNonReadable:
1174 case SpvDecorationUniform:
1175 case SpvDecorationUniformId:
1176 case SpvDecorationLocation:
1177 case SpvDecorationComponent:
1178 case SpvDecorationOffset:
1179 case SpvDecorationXfbBuffer:
1180 case SpvDecorationXfbStride:
1181 case SpvDecorationUserSemantic:
1182 vtn_warn("Decoration only allowed for struct members: %s",
1183 spirv_decoration_to_string(dec->decoration));
1184 break;
1185
1186 case SpvDecorationStream:
1187 /* We don't need to do anything here, as stream is filled up when
1188 * aplying the decoration to a variable, just check that if it is not a
1189 * struct member, it should be a struct.
1190 */
1191 vtn_assert(type->base_type == vtn_base_type_struct);
1192 break;
1193
1194 case SpvDecorationRelaxedPrecision:
1195 case SpvDecorationSpecId:
1196 case SpvDecorationInvariant:
1197 case SpvDecorationRestrict:
1198 case SpvDecorationAliased:
1199 case SpvDecorationConstant:
1200 case SpvDecorationIndex:
1201 case SpvDecorationBinding:
1202 case SpvDecorationDescriptorSet:
1203 case SpvDecorationLinkageAttributes:
1204 case SpvDecorationNoContraction:
1205 case SpvDecorationInputAttachmentIndex:
1206 vtn_warn("Decoration not allowed on types: %s",
1207 spirv_decoration_to_string(dec->decoration));
1208 break;
1209
1210 case SpvDecorationCPacked:
1211 if (b->shader->info.stage != MESA_SHADER_KERNEL)
1212 vtn_warn("Decoration only allowed for CL-style kernels: %s",
1213 spirv_decoration_to_string(dec->decoration));
1214 else
1215 type->packed = true;
1216 break;
1217
1218 case SpvDecorationSaturatedConversion:
1219 case SpvDecorationFuncParamAttr:
1220 case SpvDecorationFPRoundingMode:
1221 case SpvDecorationFPFastMathMode:
1222 case SpvDecorationAlignment:
1223 vtn_warn("Decoration only allowed for CL-style kernels: %s",
1224 spirv_decoration_to_string(dec->decoration));
1225 break;
1226
1227 case SpvDecorationUserTypeGOOGLE:
1228 /* User semantic decorations can safely be ignored by the driver. */
1229 break;
1230
1231 default:
1232 vtn_fail_with_decoration("Unhandled decoration", dec->decoration);
1233 }
1234 }
1235
1236 static unsigned
1237 translate_image_format(struct vtn_builder *b, SpvImageFormat format)
1238 {
1239 switch (format) {
1240 case SpvImageFormatUnknown: return PIPE_FORMAT_NONE;
1241 case SpvImageFormatRgba32f: return PIPE_FORMAT_R32G32B32A32_FLOAT;
1242 case SpvImageFormatRgba16f: return PIPE_FORMAT_R16G16B16A16_FLOAT;
1243 case SpvImageFormatR32f: return PIPE_FORMAT_R32_FLOAT;
1244 case SpvImageFormatRgba8: return PIPE_FORMAT_R8G8B8A8_UNORM;
1245 case SpvImageFormatRgba8Snorm: return PIPE_FORMAT_R8G8B8A8_SNORM;
1246 case SpvImageFormatRg32f: return PIPE_FORMAT_R32G32_FLOAT;
1247 case SpvImageFormatRg16f: return PIPE_FORMAT_R16G16_FLOAT;
1248 case SpvImageFormatR11fG11fB10f: return PIPE_FORMAT_R11G11B10_FLOAT;
1249 case SpvImageFormatR16f: return PIPE_FORMAT_R16_FLOAT;
1250 case SpvImageFormatRgba16: return PIPE_FORMAT_R16G16B16A16_UNORM;
1251 case SpvImageFormatRgb10A2: return PIPE_FORMAT_R10G10B10A2_UNORM;
1252 case SpvImageFormatRg16: return PIPE_FORMAT_R16G16_UNORM;
1253 case SpvImageFormatRg8: return PIPE_FORMAT_R8G8_UNORM;
1254 case SpvImageFormatR16: return PIPE_FORMAT_R16_UNORM;
1255 case SpvImageFormatR8: return PIPE_FORMAT_R8_UNORM;
1256 case SpvImageFormatRgba16Snorm: return PIPE_FORMAT_R16G16B16A16_SNORM;
1257 case SpvImageFormatRg16Snorm: return PIPE_FORMAT_R16G16_SNORM;
1258 case SpvImageFormatRg8Snorm: return PIPE_FORMAT_R8G8_SNORM;
1259 case SpvImageFormatR16Snorm: return PIPE_FORMAT_R16_SNORM;
1260 case SpvImageFormatR8Snorm: return PIPE_FORMAT_R8_SNORM;
1261 case SpvImageFormatRgba32i: return PIPE_FORMAT_R32G32B32A32_SINT;
1262 case SpvImageFormatRgba16i: return PIPE_FORMAT_R16G16B16A16_SINT;
1263 case SpvImageFormatRgba8i: return PIPE_FORMAT_R8G8B8A8_SINT;
1264 case SpvImageFormatR32i: return PIPE_FORMAT_R32_SINT;
1265 case SpvImageFormatRg32i: return PIPE_FORMAT_R32G32_SINT;
1266 case SpvImageFormatRg16i: return PIPE_FORMAT_R16G16_SINT;
1267 case SpvImageFormatRg8i: return PIPE_FORMAT_R8G8_SINT;
1268 case SpvImageFormatR16i: return PIPE_FORMAT_R16_SINT;
1269 case SpvImageFormatR8i: return PIPE_FORMAT_R8_SINT;
1270 case SpvImageFormatRgba32ui: return PIPE_FORMAT_R32G32B32A32_UINT;
1271 case SpvImageFormatRgba16ui: return PIPE_FORMAT_R16G16B16A16_UINT;
1272 case SpvImageFormatRgba8ui: return PIPE_FORMAT_R8G8B8A8_UINT;
1273 case SpvImageFormatR32ui: return PIPE_FORMAT_R32_UINT;
1274 case SpvImageFormatRgb10a2ui: return PIPE_FORMAT_R10G10B10A2_UINT;
1275 case SpvImageFormatRg32ui: return PIPE_FORMAT_R32G32_UINT;
1276 case SpvImageFormatRg16ui: return PIPE_FORMAT_R16G16_UINT;
1277 case SpvImageFormatRg8ui: return PIPE_FORMAT_R8G8_UINT;
1278 case SpvImageFormatR16ui: return PIPE_FORMAT_R16_UINT;
1279 case SpvImageFormatR8ui: return PIPE_FORMAT_R8_UINT;
1280 default:
1281 vtn_fail("Invalid image format: %s (%u)",
1282 spirv_imageformat_to_string(format), format);
1283 }
1284 }
1285
1286 static void
1287 vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
1288 const uint32_t *w, unsigned count)
1289 {
1290 struct vtn_value *val = NULL;
1291
1292 /* In order to properly handle forward declarations, we have to defer
1293 * allocation for pointer types.
1294 */
1295 if (opcode != SpvOpTypePointer && opcode != SpvOpTypeForwardPointer) {
1296 val = vtn_push_value(b, w[1], vtn_value_type_type);
1297 vtn_fail_if(val->type != NULL,
1298 "Only pointers can have forward declarations");
1299 val->type = rzalloc(b, struct vtn_type);
1300 val->type->id = w[1];
1301 }
1302
1303 switch (opcode) {
1304 case SpvOpTypeVoid:
1305 val->type->base_type = vtn_base_type_void;
1306 val->type->type = glsl_void_type();
1307 break;
1308 case SpvOpTypeBool:
1309 val->type->base_type = vtn_base_type_scalar;
1310 val->type->type = glsl_bool_type();
1311 val->type->length = 1;
1312 break;
1313 case SpvOpTypeInt: {
1314 int bit_size = w[2];
1315 const bool signedness = w[3];
1316 val->type->base_type = vtn_base_type_scalar;
1317 switch (bit_size) {
1318 case 64:
1319 val->type->type = (signedness ? glsl_int64_t_type() : glsl_uint64_t_type());
1320 break;
1321 case 32:
1322 val->type->type = (signedness ? glsl_int_type() : glsl_uint_type());
1323 break;
1324 case 16:
1325 val->type->type = (signedness ? glsl_int16_t_type() : glsl_uint16_t_type());
1326 break;
1327 case 8:
1328 val->type->type = (signedness ? glsl_int8_t_type() : glsl_uint8_t_type());
1329 break;
1330 default:
1331 vtn_fail("Invalid int bit size: %u", bit_size);
1332 }
1333 val->type->length = 1;
1334 break;
1335 }
1336
1337 case SpvOpTypeFloat: {
1338 int bit_size = w[2];
1339 val->type->base_type = vtn_base_type_scalar;
1340 switch (bit_size) {
1341 case 16:
1342 val->type->type = glsl_float16_t_type();
1343 break;
1344 case 32:
1345 val->type->type = glsl_float_type();
1346 break;
1347 case 64:
1348 val->type->type = glsl_double_type();
1349 break;
1350 default:
1351 vtn_fail("Invalid float bit size: %u", bit_size);
1352 }
1353 val->type->length = 1;
1354 break;
1355 }
1356
1357 case SpvOpTypeVector: {
1358 struct vtn_type *base = vtn_get_type(b, w[2]);
1359 unsigned elems = w[3];
1360
1361 vtn_fail_if(base->base_type != vtn_base_type_scalar,
1362 "Base type for OpTypeVector must be a scalar");
1363 vtn_fail_if((elems < 2 || elems > 4) && (elems != 8) && (elems != 16),
1364 "Invalid component count for OpTypeVector");
1365
1366 val->type->base_type = vtn_base_type_vector;
1367 val->type->type = glsl_vector_type(glsl_get_base_type(base->type), elems);
1368 val->type->length = elems;
1369 val->type->stride = glsl_type_is_boolean(val->type->type)
1370 ? 4 : glsl_get_bit_size(base->type) / 8;
1371 val->type->array_element = base;
1372 break;
1373 }
1374
1375 case SpvOpTypeMatrix: {
1376 struct vtn_type *base = vtn_get_type(b, w[2]);
1377 unsigned columns = w[3];
1378
1379 vtn_fail_if(base->base_type != vtn_base_type_vector,
1380 "Base type for OpTypeMatrix must be a vector");
1381 vtn_fail_if(columns < 2 || columns > 4,
1382 "Invalid column count for OpTypeMatrix");
1383
1384 val->type->base_type = vtn_base_type_matrix;
1385 val->type->type = glsl_matrix_type(glsl_get_base_type(base->type),
1386 glsl_get_vector_elements(base->type),
1387 columns);
1388 vtn_fail_if(glsl_type_is_error(val->type->type),
1389 "Unsupported base type for OpTypeMatrix");
1390 assert(!glsl_type_is_error(val->type->type));
1391 val->type->length = columns;
1392 val->type->array_element = base;
1393 val->type->row_major = false;
1394 val->type->stride = 0;
1395 break;
1396 }
1397
1398 case SpvOpTypeRuntimeArray:
1399 case SpvOpTypeArray: {
1400 struct vtn_type *array_element = vtn_get_type(b, w[2]);
1401
1402 if (opcode == SpvOpTypeRuntimeArray) {
1403 /* A length of 0 is used to denote unsized arrays */
1404 val->type->length = 0;
1405 } else {
1406 val->type->length = vtn_constant_uint(b, w[3]);
1407 }
1408
1409 val->type->base_type = vtn_base_type_array;
1410 val->type->array_element = array_element;
1411 if (b->shader->info.stage == MESA_SHADER_KERNEL)
1412 val->type->stride = glsl_get_cl_size(array_element->type);
1413
1414 vtn_foreach_decoration(b, val, array_stride_decoration_cb, NULL);
1415 val->type->type = glsl_array_type(array_element->type, val->type->length,
1416 val->type->stride);
1417 break;
1418 }
1419
1420 case SpvOpTypeStruct: {
1421 unsigned num_fields = count - 2;
1422 val->type->base_type = vtn_base_type_struct;
1423 val->type->length = num_fields;
1424 val->type->members = ralloc_array(b, struct vtn_type *, num_fields);
1425 val->type->offsets = ralloc_array(b, unsigned, num_fields);
1426 val->type->packed = false;
1427
1428 NIR_VLA(struct glsl_struct_field, fields, count);
1429 for (unsigned i = 0; i < num_fields; i++) {
1430 val->type->members[i] = vtn_get_type(b, w[i + 2]);
1431 fields[i] = (struct glsl_struct_field) {
1432 .type = val->type->members[i]->type,
1433 .name = ralloc_asprintf(b, "field%d", i),
1434 .location = -1,
1435 .offset = -1,
1436 };
1437 }
1438
1439 if (b->shader->info.stage == MESA_SHADER_KERNEL) {
1440 unsigned offset = 0;
1441 for (unsigned i = 0; i < num_fields; i++) {
1442 offset = align(offset, glsl_get_cl_alignment(fields[i].type));
1443 fields[i].offset = offset;
1444 offset += glsl_get_cl_size(fields[i].type);
1445 }
1446 }
1447
1448 struct member_decoration_ctx ctx = {
1449 .num_fields = num_fields,
1450 .fields = fields,
1451 .type = val->type
1452 };
1453
1454 vtn_foreach_decoration(b, val, struct_member_decoration_cb, &ctx);
1455 vtn_foreach_decoration(b, val, struct_member_matrix_stride_cb, &ctx);
1456
1457 vtn_foreach_decoration(b, val, struct_block_decoration_cb, NULL);
1458
1459 const char *name = val->name;
1460
1461 if (val->type->block || val->type->buffer_block) {
1462 /* Packing will be ignored since types coming from SPIR-V are
1463 * explicitly laid out.
1464 */
1465 val->type->type = glsl_interface_type(fields, num_fields,
1466 /* packing */ 0, false,
1467 name ? name : "block");
1468 } else {
1469 val->type->type = glsl_struct_type(fields, num_fields,
1470 name ? name : "struct", false);
1471 }
1472 break;
1473 }
1474
1475 case SpvOpTypeFunction: {
1476 val->type->base_type = vtn_base_type_function;
1477 val->type->type = NULL;
1478
1479 val->type->return_type = vtn_get_type(b, w[2]);
1480
1481 const unsigned num_params = count - 3;
1482 val->type->length = num_params;
1483 val->type->params = ralloc_array(b, struct vtn_type *, num_params);
1484 for (unsigned i = 0; i < count - 3; i++) {
1485 val->type->params[i] = vtn_get_type(b, w[i + 3]);
1486 }
1487 break;
1488 }
1489
1490 case SpvOpTypePointer:
1491 case SpvOpTypeForwardPointer: {
1492 /* We can't blindly push the value because it might be a forward
1493 * declaration.
1494 */
1495 val = vtn_untyped_value(b, w[1]);
1496
1497 SpvStorageClass storage_class = w[2];
1498
1499 if (val->value_type == vtn_value_type_invalid) {
1500 val->value_type = vtn_value_type_type;
1501 val->type = rzalloc(b, struct vtn_type);
1502 val->type->id = w[1];
1503 val->type->base_type = vtn_base_type_pointer;
1504 val->type->storage_class = storage_class;
1505
1506 /* These can actually be stored to nir_variables and used as SSA
1507 * values so they need a real glsl_type.
1508 */
1509 enum vtn_variable_mode mode = vtn_storage_class_to_mode(
1510 b, storage_class, NULL, NULL);
1511 val->type->type = nir_address_format_to_glsl_type(
1512 vtn_mode_to_address_format(b, mode));
1513 } else {
1514 vtn_fail_if(val->type->storage_class != storage_class,
1515 "The storage classes of an OpTypePointer and any "
1516 "OpTypeForwardPointers that provide forward "
1517 "declarations of it must match.");
1518 }
1519
1520 if (opcode == SpvOpTypePointer) {
1521 vtn_fail_if(val->type->deref != NULL,
1522 "While OpTypeForwardPointer can be used to provide a "
1523 "forward declaration of a pointer, OpTypePointer can "
1524 "only be used once for a given id.");
1525
1526 val->type->deref = vtn_get_type(b, w[3]);
1527
1528 /* Only certain storage classes use ArrayStride. The others (in
1529 * particular Workgroup) are expected to be laid out by the driver.
1530 */
1531 switch (storage_class) {
1532 case SpvStorageClassUniform:
1533 case SpvStorageClassPushConstant:
1534 case SpvStorageClassStorageBuffer:
1535 case SpvStorageClassPhysicalStorageBuffer:
1536 vtn_foreach_decoration(b, val, array_stride_decoration_cb, NULL);
1537 break;
1538 default:
1539 /* Nothing to do. */
1540 break;
1541 }
1542
1543 if (b->physical_ptrs) {
1544 switch (storage_class) {
1545 case SpvStorageClassFunction:
1546 case SpvStorageClassWorkgroup:
1547 case SpvStorageClassCrossWorkgroup:
1548 case SpvStorageClassUniformConstant:
1549 val->type->stride = align(glsl_get_cl_size(val->type->deref->type),
1550 glsl_get_cl_alignment(val->type->deref->type));
1551 break;
1552 default:
1553 break;
1554 }
1555 }
1556 }
1557 break;
1558 }
1559
1560 case SpvOpTypeImage: {
1561 val->type->base_type = vtn_base_type_image;
1562
1563 /* Images are represented in NIR as a scalar SSA value that is the
1564 * result of a deref instruction. An OpLoad on an OpTypeImage pointer
1565 * from UniformConstant memory just takes the NIR deref from the pointer
1566 * and turns it into an SSA value.
1567 */
1568 val->type->type = nir_address_format_to_glsl_type(
1569 vtn_mode_to_address_format(b, vtn_variable_mode_function));
1570
1571 const struct vtn_type *sampled_type = vtn_get_type(b, w[2]);
1572 if (b->shader->info.stage == MESA_SHADER_KERNEL) {
1573 vtn_fail_if(sampled_type->base_type != vtn_base_type_void,
1574 "Sampled type of OpTypeImage must be void for kernels");
1575 } else {
1576 vtn_fail_if(sampled_type->base_type != vtn_base_type_scalar ||
1577 glsl_get_bit_size(sampled_type->type) != 32,
1578 "Sampled type of OpTypeImage must be a 32-bit scalar");
1579 }
1580
1581 enum glsl_sampler_dim dim;
1582 switch ((SpvDim)w[3]) {
1583 case SpvDim1D: dim = GLSL_SAMPLER_DIM_1D; break;
1584 case SpvDim2D: dim = GLSL_SAMPLER_DIM_2D; break;
1585 case SpvDim3D: dim = GLSL_SAMPLER_DIM_3D; break;
1586 case SpvDimCube: dim = GLSL_SAMPLER_DIM_CUBE; break;
1587 case SpvDimRect: dim = GLSL_SAMPLER_DIM_RECT; break;
1588 case SpvDimBuffer: dim = GLSL_SAMPLER_DIM_BUF; break;
1589 case SpvDimSubpassData: dim = GLSL_SAMPLER_DIM_SUBPASS; break;
1590 default:
1591 vtn_fail("Invalid SPIR-V image dimensionality: %s (%u)",
1592 spirv_dim_to_string((SpvDim)w[3]), w[3]);
1593 }
1594
1595 /* w[4]: as per Vulkan spec "Validation Rules within a Module",
1596 * The “Depth” operand of OpTypeImage is ignored.
1597 */
1598 bool is_array = w[5];
1599 bool multisampled = w[6];
1600 unsigned sampled = w[7];
1601 SpvImageFormat format = w[8];
1602
1603 if (count > 9)
1604 val->type->access_qualifier = w[9];
1605 else if (b->shader->info.stage == MESA_SHADER_KERNEL)
1606 /* Per the CL C spec: If no qualifier is provided, read_only is assumed. */
1607 val->type->access_qualifier = SpvAccessQualifierReadOnly;
1608 else
1609 val->type->access_qualifier = SpvAccessQualifierReadWrite;
1610
1611 if (multisampled) {
1612 if (dim == GLSL_SAMPLER_DIM_2D)
1613 dim = GLSL_SAMPLER_DIM_MS;
1614 else if (dim == GLSL_SAMPLER_DIM_SUBPASS)
1615 dim = GLSL_SAMPLER_DIM_SUBPASS_MS;
1616 else
1617 vtn_fail("Unsupported multisampled image type");
1618 }
1619
1620 val->type->image_format = translate_image_format(b, format);
1621
1622 enum glsl_base_type sampled_base_type =
1623 glsl_get_base_type(sampled_type->type);
1624 if (sampled == 1) {
1625 val->type->glsl_image = glsl_sampler_type(dim, false, is_array,
1626 sampled_base_type);
1627 } else if (sampled == 2) {
1628 val->type->glsl_image = glsl_image_type(dim, is_array,
1629 sampled_base_type);
1630 } else if (b->shader->info.stage == MESA_SHADER_KERNEL) {
1631 val->type->glsl_image = glsl_image_type(dim, is_array,
1632 GLSL_TYPE_VOID);
1633 } else {
1634 vtn_fail("We need to know if the image will be sampled");
1635 }
1636 break;
1637 }
1638
1639 case SpvOpTypeSampledImage: {
1640 val->type->base_type = vtn_base_type_sampled_image;
1641 val->type->image = vtn_get_type(b, w[2]);
1642
1643 /* Sampled images are represented NIR as a vec2 SSA value where each
1644 * component is the result of a deref instruction. The first component
1645 * is the image and the second is the sampler. An OpLoad on an
1646 * OpTypeSampledImage pointer from UniformConstant memory just takes
1647 * the NIR deref from the pointer and duplicates it to both vector
1648 * components.
1649 */
1650 nir_address_format addr_format =
1651 vtn_mode_to_address_format(b, vtn_variable_mode_function);
1652 assert(nir_address_format_num_components(addr_format) == 1);
1653 unsigned bit_size = nir_address_format_bit_size(addr_format);
1654 assert(bit_size == 32 || bit_size == 64);
1655
1656 enum glsl_base_type base_type =
1657 bit_size == 32 ? GLSL_TYPE_UINT : GLSL_TYPE_UINT64;
1658 val->type->type = glsl_vector_type(base_type, 2);
1659 break;
1660 }
1661
1662 case SpvOpTypeSampler:
1663 val->type->base_type = vtn_base_type_sampler;
1664
1665 /* Samplers are represented in NIR as a scalar SSA value that is the
1666 * result of a deref instruction. An OpLoad on an OpTypeSampler pointer
1667 * from UniformConstant memory just takes the NIR deref from the pointer
1668 * and turns it into an SSA value.
1669 */
1670 val->type->type = nir_address_format_to_glsl_type(
1671 vtn_mode_to_address_format(b, vtn_variable_mode_function));
1672 break;
1673
1674 case SpvOpTypeOpaque:
1675 case SpvOpTypeEvent:
1676 case SpvOpTypeDeviceEvent:
1677 case SpvOpTypeReserveId:
1678 case SpvOpTypeQueue:
1679 case SpvOpTypePipe:
1680 default:
1681 vtn_fail_with_opcode("Unhandled opcode", opcode);
1682 }
1683
1684 vtn_foreach_decoration(b, val, type_decoration_cb, NULL);
1685
1686 if (val->type->base_type == vtn_base_type_struct &&
1687 (val->type->block || val->type->buffer_block)) {
1688 for (unsigned i = 0; i < val->type->length; i++) {
1689 vtn_fail_if(vtn_type_contains_block(b, val->type->members[i]),
1690 "Block and BufferBlock decorations cannot decorate a "
1691 "structure type that is nested at any level inside "
1692 "another structure type decorated with Block or "
1693 "BufferBlock.");
1694 }
1695 }
1696 }
1697
1698 static nir_constant *
1699 vtn_null_constant(struct vtn_builder *b, struct vtn_type *type)
1700 {
1701 nir_constant *c = rzalloc(b, nir_constant);
1702
1703 switch (type->base_type) {
1704 case vtn_base_type_scalar:
1705 case vtn_base_type_vector:
1706 /* Nothing to do here. It's already initialized to zero */
1707 break;
1708
1709 case vtn_base_type_pointer: {
1710 enum vtn_variable_mode mode = vtn_storage_class_to_mode(
1711 b, type->storage_class, type->deref, NULL);
1712 nir_address_format addr_format = vtn_mode_to_address_format(b, mode);
1713
1714 const nir_const_value *null_value = nir_address_format_null_value(addr_format);
1715 memcpy(c->values, null_value,
1716 sizeof(nir_const_value) * nir_address_format_num_components(addr_format));
1717 break;
1718 }
1719
1720 case vtn_base_type_void:
1721 case vtn_base_type_image:
1722 case vtn_base_type_sampler:
1723 case vtn_base_type_sampled_image:
1724 case vtn_base_type_function:
1725 /* For those we have to return something but it doesn't matter what. */
1726 break;
1727
1728 case vtn_base_type_matrix:
1729 case vtn_base_type_array:
1730 vtn_assert(type->length > 0);
1731 c->num_elements = type->length;
1732 c->elements = ralloc_array(b, nir_constant *, c->num_elements);
1733
1734 c->elements[0] = vtn_null_constant(b, type->array_element);
1735 for (unsigned i = 1; i < c->num_elements; i++)
1736 c->elements[i] = c->elements[0];
1737 break;
1738
1739 case vtn_base_type_struct:
1740 c->num_elements = type->length;
1741 c->elements = ralloc_array(b, nir_constant *, c->num_elements);
1742 for (unsigned i = 0; i < c->num_elements; i++)
1743 c->elements[i] = vtn_null_constant(b, type->members[i]);
1744 break;
1745
1746 default:
1747 vtn_fail("Invalid type for null constant");
1748 }
1749
1750 return c;
1751 }
1752
1753 static void
1754 spec_constant_decoration_cb(struct vtn_builder *b, UNUSED struct vtn_value *val,
1755 ASSERTED int member,
1756 const struct vtn_decoration *dec, void *data)
1757 {
1758 vtn_assert(member == -1);
1759 if (dec->decoration != SpvDecorationSpecId)
1760 return;
1761
1762 nir_const_value *value = data;
1763 for (unsigned i = 0; i < b->num_specializations; i++) {
1764 if (b->specializations[i].id == dec->operands[0]) {
1765 *value = b->specializations[i].value;
1766 return;
1767 }
1768 }
1769 }
1770
1771 static void
1772 handle_workgroup_size_decoration_cb(struct vtn_builder *b,
1773 struct vtn_value *val,
1774 ASSERTED int member,
1775 const struct vtn_decoration *dec,
1776 UNUSED void *data)
1777 {
1778 vtn_assert(member == -1);
1779 if (dec->decoration != SpvDecorationBuiltIn ||
1780 dec->operands[0] != SpvBuiltInWorkgroupSize)
1781 return;
1782
1783 vtn_assert(val->type->type == glsl_vector_type(GLSL_TYPE_UINT, 3));
1784 b->workgroup_size_builtin = val;
1785 }
1786
1787 static void
1788 vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
1789 const uint32_t *w, unsigned count)
1790 {
1791 struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_constant);
1792 val->constant = rzalloc(b, nir_constant);
1793 switch (opcode) {
1794 case SpvOpConstantTrue:
1795 case SpvOpConstantFalse:
1796 case SpvOpSpecConstantTrue:
1797 case SpvOpSpecConstantFalse: {
1798 vtn_fail_if(val->type->type != glsl_bool_type(),
1799 "Result type of %s must be OpTypeBool",
1800 spirv_op_to_string(opcode));
1801
1802 bool bval = (opcode == SpvOpConstantTrue ||
1803 opcode == SpvOpSpecConstantTrue);
1804
1805 nir_const_value u32val = nir_const_value_for_uint(bval, 32);
1806
1807 if (opcode == SpvOpSpecConstantTrue ||
1808 opcode == SpvOpSpecConstantFalse)
1809 vtn_foreach_decoration(b, val, spec_constant_decoration_cb, &u32val);
1810
1811 val->constant->values[0].b = u32val.u32 != 0;
1812 break;
1813 }
1814
1815 case SpvOpConstant:
1816 case SpvOpSpecConstant: {
1817 vtn_fail_if(val->type->base_type != vtn_base_type_scalar,
1818 "Result type of %s must be a scalar",
1819 spirv_op_to_string(opcode));
1820 int bit_size = glsl_get_bit_size(val->type->type);
1821 switch (bit_size) {
1822 case 64:
1823 val->constant->values[0].u64 = vtn_u64_literal(&w[3]);
1824 break;
1825 case 32:
1826 val->constant->values[0].u32 = w[3];
1827 break;
1828 case 16:
1829 val->constant->values[0].u16 = w[3];
1830 break;
1831 case 8:
1832 val->constant->values[0].u8 = w[3];
1833 break;
1834 default:
1835 vtn_fail("Unsupported SpvOpConstant bit size: %u", bit_size);
1836 }
1837
1838 if (opcode == SpvOpSpecConstant)
1839 vtn_foreach_decoration(b, val, spec_constant_decoration_cb,
1840 &val->constant->values[0]);
1841 break;
1842 }
1843
1844 case SpvOpSpecConstantComposite:
1845 case SpvOpConstantComposite: {
1846 unsigned elem_count = count - 3;
1847 vtn_fail_if(elem_count != val->type->length,
1848 "%s has %u constituents, expected %u",
1849 spirv_op_to_string(opcode), elem_count, val->type->length);
1850
1851 nir_constant **elems = ralloc_array(b, nir_constant *, elem_count);
1852 for (unsigned i = 0; i < elem_count; i++) {
1853 struct vtn_value *val = vtn_untyped_value(b, w[i + 3]);
1854
1855 if (val->value_type == vtn_value_type_constant) {
1856 elems[i] = val->constant;
1857 } else {
1858 vtn_fail_if(val->value_type != vtn_value_type_undef,
1859 "only constants or undefs allowed for "
1860 "SpvOpConstantComposite");
1861 /* to make it easier, just insert a NULL constant for now */
1862 elems[i] = vtn_null_constant(b, val->type);
1863 }
1864 }
1865
1866 switch (val->type->base_type) {
1867 case vtn_base_type_vector: {
1868 assert(glsl_type_is_vector(val->type->type));
1869 for (unsigned i = 0; i < elem_count; i++)
1870 val->constant->values[i] = elems[i]->values[0];
1871 break;
1872 }
1873
1874 case vtn_base_type_matrix:
1875 case vtn_base_type_struct:
1876 case vtn_base_type_array:
1877 ralloc_steal(val->constant, elems);
1878 val->constant->num_elements = elem_count;
1879 val->constant->elements = elems;
1880 break;
1881
1882 default:
1883 vtn_fail("Result type of %s must be a composite type",
1884 spirv_op_to_string(opcode));
1885 }
1886 break;
1887 }
1888
1889 case SpvOpSpecConstantOp: {
1890 nir_const_value u32op = nir_const_value_for_uint(w[3], 32);
1891 vtn_foreach_decoration(b, val, spec_constant_decoration_cb, &u32op);
1892 SpvOp opcode = u32op.u32;
1893 switch (opcode) {
1894 case SpvOpVectorShuffle: {
1895 struct vtn_value *v0 = &b->values[w[4]];
1896 struct vtn_value *v1 = &b->values[w[5]];
1897
1898 vtn_assert(v0->value_type == vtn_value_type_constant ||
1899 v0->value_type == vtn_value_type_undef);
1900 vtn_assert(v1->value_type == vtn_value_type_constant ||
1901 v1->value_type == vtn_value_type_undef);
1902
1903 unsigned len0 = glsl_get_vector_elements(v0->type->type);
1904 unsigned len1 = glsl_get_vector_elements(v1->type->type);
1905
1906 vtn_assert(len0 + len1 < 16);
1907
1908 unsigned bit_size = glsl_get_bit_size(val->type->type);
1909 unsigned bit_size0 = glsl_get_bit_size(v0->type->type);
1910 unsigned bit_size1 = glsl_get_bit_size(v1->type->type);
1911
1912 vtn_assert(bit_size == bit_size0 && bit_size == bit_size1);
1913 (void)bit_size0; (void)bit_size1;
1914
1915 nir_const_value undef = { .u64 = 0xdeadbeefdeadbeef };
1916 nir_const_value combined[NIR_MAX_VEC_COMPONENTS * 2];
1917
1918 if (v0->value_type == vtn_value_type_constant) {
1919 for (unsigned i = 0; i < len0; i++)
1920 combined[i] = v0->constant->values[i];
1921 }
1922 if (v1->value_type == vtn_value_type_constant) {
1923 for (unsigned i = 0; i < len1; i++)
1924 combined[len0 + i] = v1->constant->values[i];
1925 }
1926
1927 for (unsigned i = 0, j = 0; i < count - 6; i++, j++) {
1928 uint32_t comp = w[i + 6];
1929 if (comp == (uint32_t)-1) {
1930 /* If component is not used, set the value to a known constant
1931 * to detect if it is wrongly used.
1932 */
1933 val->constant->values[j] = undef;
1934 } else {
1935 vtn_fail_if(comp >= len0 + len1,
1936 "All Component literals must either be FFFFFFFF "
1937 "or in [0, N - 1] (inclusive).");
1938 val->constant->values[j] = combined[comp];
1939 }
1940 }
1941 break;
1942 }
1943
1944 case SpvOpCompositeExtract:
1945 case SpvOpCompositeInsert: {
1946 struct vtn_value *comp;
1947 unsigned deref_start;
1948 struct nir_constant **c;
1949 if (opcode == SpvOpCompositeExtract) {
1950 comp = vtn_value(b, w[4], vtn_value_type_constant);
1951 deref_start = 5;
1952 c = &comp->constant;
1953 } else {
1954 comp = vtn_value(b, w[5], vtn_value_type_constant);
1955 deref_start = 6;
1956 val->constant = nir_constant_clone(comp->constant,
1957 (nir_variable *)b);
1958 c = &val->constant;
1959 }
1960
1961 int elem = -1;
1962 const struct vtn_type *type = comp->type;
1963 for (unsigned i = deref_start; i < count; i++) {
1964 vtn_fail_if(w[i] > type->length,
1965 "%uth index of %s is %u but the type has only "
1966 "%u elements", i - deref_start,
1967 spirv_op_to_string(opcode), w[i], type->length);
1968
1969 switch (type->base_type) {
1970 case vtn_base_type_vector:
1971 elem = w[i];
1972 type = type->array_element;
1973 break;
1974
1975 case vtn_base_type_matrix:
1976 case vtn_base_type_array:
1977 c = &(*c)->elements[w[i]];
1978 type = type->array_element;
1979 break;
1980
1981 case vtn_base_type_struct:
1982 c = &(*c)->elements[w[i]];
1983 type = type->members[w[i]];
1984 break;
1985
1986 default:
1987 vtn_fail("%s must only index into composite types",
1988 spirv_op_to_string(opcode));
1989 }
1990 }
1991
1992 if (opcode == SpvOpCompositeExtract) {
1993 if (elem == -1) {
1994 val->constant = *c;
1995 } else {
1996 unsigned num_components = type->length;
1997 for (unsigned i = 0; i < num_components; i++)
1998 val->constant->values[i] = (*c)->values[elem + i];
1999 }
2000 } else {
2001 struct vtn_value *insert =
2002 vtn_value(b, w[4], vtn_value_type_constant);
2003 vtn_assert(insert->type == type);
2004 if (elem == -1) {
2005 *c = insert->constant;
2006 } else {
2007 unsigned num_components = type->length;
2008 for (unsigned i = 0; i < num_components; i++)
2009 (*c)->values[elem + i] = insert->constant->values[i];
2010 }
2011 }
2012 break;
2013 }
2014
2015 default: {
2016 bool swap;
2017 nir_alu_type dst_alu_type = nir_get_nir_type_for_glsl_type(val->type->type);
2018 nir_alu_type src_alu_type = dst_alu_type;
2019 unsigned num_components = glsl_get_vector_elements(val->type->type);
2020 unsigned bit_size;
2021
2022 vtn_assert(count <= 7);
2023
2024 switch (opcode) {
2025 case SpvOpSConvert:
2026 case SpvOpFConvert:
2027 case SpvOpUConvert:
2028 /* We have a source in a conversion */
2029 src_alu_type =
2030 nir_get_nir_type_for_glsl_type(vtn_get_value_type(b, w[4])->type);
2031 /* We use the bitsize of the conversion source to evaluate the opcode later */
2032 bit_size = glsl_get_bit_size(vtn_get_value_type(b, w[4])->type);
2033 break;
2034 default:
2035 bit_size = glsl_get_bit_size(val->type->type);
2036 };
2037
2038 nir_op op = vtn_nir_alu_op_for_spirv_opcode(b, opcode, &swap,
2039 nir_alu_type_get_type_size(src_alu_type),
2040 nir_alu_type_get_type_size(dst_alu_type));
2041 nir_const_value src[3][NIR_MAX_VEC_COMPONENTS];
2042
2043 for (unsigned i = 0; i < count - 4; i++) {
2044 struct vtn_value *src_val =
2045 vtn_value(b, w[4 + i], vtn_value_type_constant);
2046
2047 /* If this is an unsized source, pull the bit size from the
2048 * source; otherwise, we'll use the bit size from the destination.
2049 */
2050 if (!nir_alu_type_get_type_size(nir_op_infos[op].input_types[i]))
2051 bit_size = glsl_get_bit_size(src_val->type->type);
2052
2053 unsigned src_comps = nir_op_infos[op].input_sizes[i] ?
2054 nir_op_infos[op].input_sizes[i] :
2055 num_components;
2056
2057 unsigned j = swap ? 1 - i : i;
2058 for (unsigned c = 0; c < src_comps; c++)
2059 src[j][c] = src_val->constant->values[c];
2060 }
2061
2062 /* fix up fixed size sources */
2063 switch (op) {
2064 case nir_op_ishl:
2065 case nir_op_ishr:
2066 case nir_op_ushr: {
2067 if (bit_size == 32)
2068 break;
2069 for (unsigned i = 0; i < num_components; ++i) {
2070 switch (bit_size) {
2071 case 64: src[1][i].u32 = src[1][i].u64; break;
2072 case 16: src[1][i].u32 = src[1][i].u16; break;
2073 case 8: src[1][i].u32 = src[1][i].u8; break;
2074 }
2075 }
2076 break;
2077 }
2078 default:
2079 break;
2080 }
2081
2082 nir_const_value *srcs[3] = {
2083 src[0], src[1], src[2],
2084 };
2085 nir_eval_const_opcode(op, val->constant->values,
2086 num_components, bit_size, srcs,
2087 b->shader->info.float_controls_execution_mode);
2088 break;
2089 } /* default */
2090 }
2091 break;
2092 }
2093
2094 case SpvOpConstantNull:
2095 val->constant = vtn_null_constant(b, val->type);
2096 break;
2097
2098 default:
2099 vtn_fail_with_opcode("Unhandled opcode", opcode);
2100 }
2101
2102 /* Now that we have the value, update the workgroup size if needed */
2103 vtn_foreach_decoration(b, val, handle_workgroup_size_decoration_cb, NULL);
2104 }
2105
2106 static void
2107 vtn_split_barrier_semantics(struct vtn_builder *b,
2108 SpvMemorySemanticsMask semantics,
2109 SpvMemorySemanticsMask *before,
2110 SpvMemorySemanticsMask *after)
2111 {
2112 /* For memory semantics embedded in operations, we split them into up to
2113 * two barriers, to be added before and after the operation. This is less
2114 * strict than if we propagated until the final backend stage, but still
2115 * result in correct execution.
2116 *
2117 * A further improvement could be pipe this information (and use!) into the
2118 * next compiler layers, at the expense of making the handling of barriers
2119 * more complicated.
2120 */
2121
2122 *before = SpvMemorySemanticsMaskNone;
2123 *after = SpvMemorySemanticsMaskNone;
2124
2125 SpvMemorySemanticsMask order_semantics =
2126 semantics & (SpvMemorySemanticsAcquireMask |
2127 SpvMemorySemanticsReleaseMask |
2128 SpvMemorySemanticsAcquireReleaseMask |
2129 SpvMemorySemanticsSequentiallyConsistentMask);
2130
2131 if (util_bitcount(order_semantics) > 1) {
2132 /* Old GLSLang versions incorrectly set all the ordering bits. This was
2133 * fixed in c51287d744fb6e7e9ccc09f6f8451e6c64b1dad6 of glslang repo,
2134 * and it is in GLSLang since revision "SPIRV99.1321" (from Jul-2016).
2135 */
2136 vtn_warn("Multiple memory ordering semantics specified, "
2137 "assuming AcquireRelease.");
2138 order_semantics = SpvMemorySemanticsAcquireReleaseMask;
2139 }
2140
2141 const SpvMemorySemanticsMask av_vis_semantics =
2142 semantics & (SpvMemorySemanticsMakeAvailableMask |
2143 SpvMemorySemanticsMakeVisibleMask);
2144
2145 const SpvMemorySemanticsMask storage_semantics =
2146 semantics & (SpvMemorySemanticsUniformMemoryMask |
2147 SpvMemorySemanticsSubgroupMemoryMask |
2148 SpvMemorySemanticsWorkgroupMemoryMask |
2149 SpvMemorySemanticsCrossWorkgroupMemoryMask |
2150 SpvMemorySemanticsAtomicCounterMemoryMask |
2151 SpvMemorySemanticsImageMemoryMask |
2152 SpvMemorySemanticsOutputMemoryMask);
2153
2154 const SpvMemorySemanticsMask other_semantics =
2155 semantics & ~(order_semantics | av_vis_semantics | storage_semantics |
2156 SpvMemorySemanticsVolatileMask);
2157
2158 if (other_semantics)
2159 vtn_warn("Ignoring unhandled memory semantics: %u\n", other_semantics);
2160
2161 /* SequentiallyConsistent is treated as AcquireRelease. */
2162
2163 /* The RELEASE barrier happens BEFORE the operation, and it is usually
2164 * associated with a Store. All the write operations with a matching
2165 * semantics will not be reordered after the Store.
2166 */
2167 if (order_semantics & (SpvMemorySemanticsReleaseMask |
2168 SpvMemorySemanticsAcquireReleaseMask |
2169 SpvMemorySemanticsSequentiallyConsistentMask)) {
2170 *before |= SpvMemorySemanticsReleaseMask | storage_semantics;
2171 }
2172
2173 /* The ACQUIRE barrier happens AFTER the operation, and it is usually
2174 * associated with a Load. All the operations with a matching semantics
2175 * will not be reordered before the Load.
2176 */
2177 if (order_semantics & (SpvMemorySemanticsAcquireMask |
2178 SpvMemorySemanticsAcquireReleaseMask |
2179 SpvMemorySemanticsSequentiallyConsistentMask)) {
2180 *after |= SpvMemorySemanticsAcquireMask | storage_semantics;
2181 }
2182
2183 if (av_vis_semantics & SpvMemorySemanticsMakeVisibleMask)
2184 *before |= SpvMemorySemanticsMakeVisibleMask | storage_semantics;
2185
2186 if (av_vis_semantics & SpvMemorySemanticsMakeAvailableMask)
2187 *after |= SpvMemorySemanticsMakeAvailableMask | storage_semantics;
2188 }
2189
2190 static nir_memory_semantics
2191 vtn_mem_semantics_to_nir_mem_semantics(struct vtn_builder *b,
2192 SpvMemorySemanticsMask semantics)
2193 {
2194 nir_memory_semantics nir_semantics = 0;
2195
2196 SpvMemorySemanticsMask order_semantics =
2197 semantics & (SpvMemorySemanticsAcquireMask |
2198 SpvMemorySemanticsReleaseMask |
2199 SpvMemorySemanticsAcquireReleaseMask |
2200 SpvMemorySemanticsSequentiallyConsistentMask);
2201
2202 if (util_bitcount(order_semantics) > 1) {
2203 /* Old GLSLang versions incorrectly set all the ordering bits. This was
2204 * fixed in c51287d744fb6e7e9ccc09f6f8451e6c64b1dad6 of glslang repo,
2205 * and it is in GLSLang since revision "SPIRV99.1321" (from Jul-2016).
2206 */
2207 vtn_warn("Multiple memory ordering semantics bits specified, "
2208 "assuming AcquireRelease.");
2209 order_semantics = SpvMemorySemanticsAcquireReleaseMask;
2210 }
2211
2212 switch (order_semantics) {
2213 case 0:
2214 /* Not an ordering barrier. */
2215 break;
2216
2217 case SpvMemorySemanticsAcquireMask:
2218 nir_semantics = NIR_MEMORY_ACQUIRE;
2219 break;
2220
2221 case SpvMemorySemanticsReleaseMask:
2222 nir_semantics = NIR_MEMORY_RELEASE;
2223 break;
2224
2225 case SpvMemorySemanticsSequentiallyConsistentMask:
2226 /* Fall through. Treated as AcquireRelease in Vulkan. */
2227 case SpvMemorySemanticsAcquireReleaseMask:
2228 nir_semantics = NIR_MEMORY_ACQUIRE | NIR_MEMORY_RELEASE;
2229 break;
2230
2231 default:
2232 unreachable("Invalid memory order semantics");
2233 }
2234
2235 if (semantics & SpvMemorySemanticsMakeAvailableMask) {
2236 vtn_fail_if(!b->options->caps.vk_memory_model,
2237 "To use MakeAvailable memory semantics the VulkanMemoryModel "
2238 "capability must be declared.");
2239 nir_semantics |= NIR_MEMORY_MAKE_AVAILABLE;
2240 }
2241
2242 if (semantics & SpvMemorySemanticsMakeVisibleMask) {
2243 vtn_fail_if(!b->options->caps.vk_memory_model,
2244 "To use MakeVisible memory semantics the VulkanMemoryModel "
2245 "capability must be declared.");
2246 nir_semantics |= NIR_MEMORY_MAKE_VISIBLE;
2247 }
2248
2249 return nir_semantics;
2250 }
2251
2252 static nir_variable_mode
2253 vtn_mem_sematics_to_nir_var_modes(struct vtn_builder *b,
2254 SpvMemorySemanticsMask semantics)
2255 {
2256 /* Vulkan Environment for SPIR-V says "SubgroupMemory, CrossWorkgroupMemory,
2257 * and AtomicCounterMemory are ignored".
2258 */
2259 semantics &= ~(SpvMemorySemanticsSubgroupMemoryMask |
2260 SpvMemorySemanticsCrossWorkgroupMemoryMask |
2261 SpvMemorySemanticsAtomicCounterMemoryMask);
2262
2263 /* TODO: Consider adding nir_var_mem_image mode to NIR so it can be used
2264 * for SpvMemorySemanticsImageMemoryMask.
2265 */
2266
2267 nir_variable_mode modes = 0;
2268 if (semantics & (SpvMemorySemanticsUniformMemoryMask |
2269 SpvMemorySemanticsImageMemoryMask)) {
2270 modes |= nir_var_uniform |
2271 nir_var_mem_ubo |
2272 nir_var_mem_ssbo |
2273 nir_var_mem_global;
2274 }
2275 if (semantics & SpvMemorySemanticsWorkgroupMemoryMask)
2276 modes |= nir_var_mem_shared;
2277 if (semantics & SpvMemorySemanticsOutputMemoryMask) {
2278 modes |= nir_var_shader_out;
2279 }
2280
2281 return modes;
2282 }
2283
2284 static nir_scope
2285 vtn_scope_to_nir_scope(struct vtn_builder *b, SpvScope scope)
2286 {
2287 nir_scope nir_scope;
2288 switch (scope) {
2289 case SpvScopeDevice:
2290 vtn_fail_if(b->options->caps.vk_memory_model &&
2291 !b->options->caps.vk_memory_model_device_scope,
2292 "If the Vulkan memory model is declared and any instruction "
2293 "uses Device scope, the VulkanMemoryModelDeviceScope "
2294 "capability must be declared.");
2295 nir_scope = NIR_SCOPE_DEVICE;
2296 break;
2297
2298 case SpvScopeQueueFamily:
2299 vtn_fail_if(!b->options->caps.vk_memory_model,
2300 "To use Queue Family scope, the VulkanMemoryModel capability "
2301 "must be declared.");
2302 nir_scope = NIR_SCOPE_QUEUE_FAMILY;
2303 break;
2304
2305 case SpvScopeWorkgroup:
2306 nir_scope = NIR_SCOPE_WORKGROUP;
2307 break;
2308
2309 case SpvScopeSubgroup:
2310 nir_scope = NIR_SCOPE_SUBGROUP;
2311 break;
2312
2313 case SpvScopeInvocation:
2314 nir_scope = NIR_SCOPE_INVOCATION;
2315 break;
2316
2317 default:
2318 vtn_fail("Invalid memory scope");
2319 }
2320
2321 return nir_scope;
2322 }
2323
2324 static void
2325 vtn_emit_scoped_control_barrier(struct vtn_builder *b, SpvScope exec_scope,
2326 SpvScope mem_scope,
2327 SpvMemorySemanticsMask semantics)
2328 {
2329 nir_memory_semantics nir_semantics =
2330 vtn_mem_semantics_to_nir_mem_semantics(b, semantics);
2331 nir_variable_mode modes = vtn_mem_sematics_to_nir_var_modes(b, semantics);
2332 nir_scope nir_exec_scope = vtn_scope_to_nir_scope(b, exec_scope);
2333
2334 /* Memory semantics is optional for OpControlBarrier. */
2335 nir_scope nir_mem_scope;
2336 if (nir_semantics == 0 || modes == 0)
2337 nir_mem_scope = NIR_SCOPE_NONE;
2338 else
2339 nir_mem_scope = vtn_scope_to_nir_scope(b, mem_scope);
2340
2341 nir_scoped_barrier(&b->nb, nir_exec_scope, nir_mem_scope, nir_semantics, modes);
2342 }
2343
2344 static void
2345 vtn_emit_scoped_memory_barrier(struct vtn_builder *b, SpvScope scope,
2346 SpvMemorySemanticsMask semantics)
2347 {
2348 nir_variable_mode modes = vtn_mem_sematics_to_nir_var_modes(b, semantics);
2349 nir_memory_semantics nir_semantics =
2350 vtn_mem_semantics_to_nir_mem_semantics(b, semantics);
2351
2352 /* No barrier to add. */
2353 if (nir_semantics == 0 || modes == 0)
2354 return;
2355
2356 nir_scope nir_mem_scope = vtn_scope_to_nir_scope(b, scope);
2357 nir_scoped_barrier(&b->nb, NIR_SCOPE_NONE, nir_mem_scope, nir_semantics, modes);
2358 }
2359
2360 struct vtn_ssa_value *
2361 vtn_create_ssa_value(struct vtn_builder *b, const struct glsl_type *type)
2362 {
2363 /* Always use bare types for SSA values for a couple of reasons:
2364 *
2365 * 1. Code which emits deref chains should never listen to the explicit
2366 * layout information on the SSA value if any exists. If we've
2367 * accidentally been relying on this, we want to find those bugs.
2368 *
2369 * 2. We want to be able to quickly check that an SSA value being assigned
2370 * to a SPIR-V value has the right type. Using bare types everywhere
2371 * ensures that we can pointer-compare.
2372 */
2373 struct vtn_ssa_value *val = rzalloc(b, struct vtn_ssa_value);
2374 val->type = glsl_get_bare_type(type);
2375
2376
2377 if (!glsl_type_is_vector_or_scalar(type)) {
2378 unsigned elems = glsl_get_length(val->type);
2379 val->elems = ralloc_array(b, struct vtn_ssa_value *, elems);
2380 if (glsl_type_is_array_or_matrix(type)) {
2381 const struct glsl_type *elem_type = glsl_get_array_element(type);
2382 for (unsigned i = 0; i < elems; i++)
2383 val->elems[i] = vtn_create_ssa_value(b, elem_type);
2384 } else {
2385 vtn_assert(glsl_type_is_struct_or_ifc(type));
2386 for (unsigned i = 0; i < elems; i++) {
2387 const struct glsl_type *elem_type = glsl_get_struct_field(type, i);
2388 val->elems[i] = vtn_create_ssa_value(b, elem_type);
2389 }
2390 }
2391 }
2392
2393 return val;
2394 }
2395
2396 static nir_tex_src
2397 vtn_tex_src(struct vtn_builder *b, unsigned index, nir_tex_src_type type)
2398 {
2399 nir_tex_src src;
2400 src.src = nir_src_for_ssa(vtn_get_nir_ssa(b, index));
2401 src.src_type = type;
2402 return src;
2403 }
2404
2405 static uint32_t
2406 image_operand_arg(struct vtn_builder *b, const uint32_t *w, uint32_t count,
2407 uint32_t mask_idx, SpvImageOperandsMask op)
2408 {
2409 static const SpvImageOperandsMask ops_with_arg =
2410 SpvImageOperandsBiasMask |
2411 SpvImageOperandsLodMask |
2412 SpvImageOperandsGradMask |
2413 SpvImageOperandsConstOffsetMask |
2414 SpvImageOperandsOffsetMask |
2415 SpvImageOperandsConstOffsetsMask |
2416 SpvImageOperandsSampleMask |
2417 SpvImageOperandsMinLodMask |
2418 SpvImageOperandsMakeTexelAvailableMask |
2419 SpvImageOperandsMakeTexelVisibleMask;
2420
2421 assert(util_bitcount(op) == 1);
2422 assert(w[mask_idx] & op);
2423 assert(op & ops_with_arg);
2424
2425 uint32_t idx = util_bitcount(w[mask_idx] & (op - 1) & ops_with_arg) + 1;
2426
2427 /* Adjust indices for operands with two arguments. */
2428 static const SpvImageOperandsMask ops_with_two_args =
2429 SpvImageOperandsGradMask;
2430 idx += util_bitcount(w[mask_idx] & (op - 1) & ops_with_two_args);
2431
2432 idx += mask_idx;
2433
2434 vtn_fail_if(idx + (op & ops_with_two_args ? 1 : 0) >= count,
2435 "Image op claims to have %s but does not enough "
2436 "following operands", spirv_imageoperands_to_string(op));
2437
2438 return idx;
2439 }
2440
2441 static void
2442 non_uniform_decoration_cb(struct vtn_builder *b,
2443 struct vtn_value *val, int member,
2444 const struct vtn_decoration *dec, void *void_ctx)
2445 {
2446 enum gl_access_qualifier *access = void_ctx;
2447 switch (dec->decoration) {
2448 case SpvDecorationNonUniformEXT:
2449 *access |= ACCESS_NON_UNIFORM;
2450 break;
2451
2452 default:
2453 break;
2454 }
2455 }
2456
2457 static void
2458 vtn_handle_texture(struct vtn_builder *b, SpvOp opcode,
2459 const uint32_t *w, unsigned count)
2460 {
2461 struct vtn_type *ret_type = vtn_get_type(b, w[1]);
2462
2463 if (opcode == SpvOpSampledImage) {
2464 struct vtn_sampled_image si = {
2465 .image = vtn_get_image(b, w[3]),
2466 .sampler = vtn_get_sampler(b, w[4]),
2467 };
2468 vtn_push_sampled_image(b, w[2], si);
2469 return;
2470 } else if (opcode == SpvOpImage) {
2471 struct vtn_sampled_image si = vtn_get_sampled_image(b, w[3]);
2472 vtn_push_image(b, w[2], si.image);
2473 return;
2474 }
2475
2476 nir_deref_instr *image = NULL, *sampler = NULL;
2477 struct vtn_value *sampled_val = vtn_untyped_value(b, w[3]);
2478 if (sampled_val->type->base_type == vtn_base_type_sampled_image) {
2479 struct vtn_sampled_image si = vtn_get_sampled_image(b, w[3]);
2480 image = si.image;
2481 sampler = si.sampler;
2482 } else {
2483 image = vtn_get_image(b, w[3]);
2484 }
2485
2486 const enum glsl_sampler_dim sampler_dim = glsl_get_sampler_dim(image->type);
2487 const bool is_array = glsl_sampler_type_is_array(image->type);
2488 nir_alu_type dest_type = nir_type_invalid;
2489
2490 /* Figure out the base texture operation */
2491 nir_texop texop;
2492 switch (opcode) {
2493 case SpvOpImageSampleImplicitLod:
2494 case SpvOpImageSampleDrefImplicitLod:
2495 case SpvOpImageSampleProjImplicitLod:
2496 case SpvOpImageSampleProjDrefImplicitLod:
2497 texop = nir_texop_tex;
2498 break;
2499
2500 case SpvOpImageSampleExplicitLod:
2501 case SpvOpImageSampleDrefExplicitLod:
2502 case SpvOpImageSampleProjExplicitLod:
2503 case SpvOpImageSampleProjDrefExplicitLod:
2504 texop = nir_texop_txl;
2505 break;
2506
2507 case SpvOpImageFetch:
2508 if (sampler_dim == GLSL_SAMPLER_DIM_MS) {
2509 texop = nir_texop_txf_ms;
2510 } else {
2511 texop = nir_texop_txf;
2512 }
2513 break;
2514
2515 case SpvOpImageGather:
2516 case SpvOpImageDrefGather:
2517 texop = nir_texop_tg4;
2518 break;
2519
2520 case SpvOpImageQuerySizeLod:
2521 case SpvOpImageQuerySize:
2522 texop = nir_texop_txs;
2523 dest_type = nir_type_int;
2524 break;
2525
2526 case SpvOpImageQueryLod:
2527 texop = nir_texop_lod;
2528 dest_type = nir_type_float;
2529 break;
2530
2531 case SpvOpImageQueryLevels:
2532 texop = nir_texop_query_levels;
2533 dest_type = nir_type_int;
2534 break;
2535
2536 case SpvOpImageQuerySamples:
2537 texop = nir_texop_texture_samples;
2538 dest_type = nir_type_int;
2539 break;
2540
2541 case SpvOpFragmentFetchAMD:
2542 texop = nir_texop_fragment_fetch;
2543 break;
2544
2545 case SpvOpFragmentMaskFetchAMD:
2546 texop = nir_texop_fragment_mask_fetch;
2547 break;
2548
2549 default:
2550 vtn_fail_with_opcode("Unhandled opcode", opcode);
2551 }
2552
2553 nir_tex_src srcs[10]; /* 10 should be enough */
2554 nir_tex_src *p = srcs;
2555
2556 p->src = nir_src_for_ssa(&image->dest.ssa);
2557 p->src_type = nir_tex_src_texture_deref;
2558 p++;
2559
2560 switch (texop) {
2561 case nir_texop_tex:
2562 case nir_texop_txb:
2563 case nir_texop_txl:
2564 case nir_texop_txd:
2565 case nir_texop_tg4:
2566 case nir_texop_lod:
2567 vtn_fail_if(sampler == NULL,
2568 "%s requires an image of type OpTypeSampledImage",
2569 spirv_op_to_string(opcode));
2570 p->src = nir_src_for_ssa(&sampler->dest.ssa);
2571 p->src_type = nir_tex_src_sampler_deref;
2572 p++;
2573 break;
2574 case nir_texop_txf:
2575 case nir_texop_txf_ms:
2576 case nir_texop_txs:
2577 case nir_texop_query_levels:
2578 case nir_texop_texture_samples:
2579 case nir_texop_samples_identical:
2580 case nir_texop_fragment_fetch:
2581 case nir_texop_fragment_mask_fetch:
2582 /* These don't */
2583 break;
2584 case nir_texop_txf_ms_fb:
2585 vtn_fail("unexpected nir_texop_txf_ms_fb");
2586 break;
2587 case nir_texop_txf_ms_mcs:
2588 vtn_fail("unexpected nir_texop_txf_ms_mcs");
2589 case nir_texop_tex_prefetch:
2590 vtn_fail("unexpected nir_texop_tex_prefetch");
2591 }
2592
2593 unsigned idx = 4;
2594
2595 struct nir_ssa_def *coord;
2596 unsigned coord_components;
2597 switch (opcode) {
2598 case SpvOpImageSampleImplicitLod:
2599 case SpvOpImageSampleExplicitLod:
2600 case SpvOpImageSampleDrefImplicitLod:
2601 case SpvOpImageSampleDrefExplicitLod:
2602 case SpvOpImageSampleProjImplicitLod:
2603 case SpvOpImageSampleProjExplicitLod:
2604 case SpvOpImageSampleProjDrefImplicitLod:
2605 case SpvOpImageSampleProjDrefExplicitLod:
2606 case SpvOpImageFetch:
2607 case SpvOpImageGather:
2608 case SpvOpImageDrefGather:
2609 case SpvOpImageQueryLod:
2610 case SpvOpFragmentFetchAMD:
2611 case SpvOpFragmentMaskFetchAMD: {
2612 /* All these types have the coordinate as their first real argument */
2613 coord_components = glsl_get_sampler_dim_coordinate_components(sampler_dim);
2614
2615 if (is_array && texop != nir_texop_lod)
2616 coord_components++;
2617
2618 struct vtn_ssa_value *coord_val = vtn_ssa_value(b, w[idx++]);
2619 coord = coord_val->def;
2620 p->src = nir_src_for_ssa(nir_channels(&b->nb, coord,
2621 (1 << coord_components) - 1));
2622
2623 /* OpenCL allows integer sampling coordinates */
2624 if (glsl_type_is_integer(coord_val->type) &&
2625 opcode == SpvOpImageSampleExplicitLod) {
2626 vtn_fail_if(b->shader->info.stage != MESA_SHADER_KERNEL,
2627 "Unless the Kernel capability is being used, the coordinate parameter "
2628 "OpImageSampleExplicitLod must be floating point.");
2629
2630 p->src = nir_src_for_ssa(nir_i2f32(&b->nb, p->src.ssa));
2631 }
2632
2633 p->src_type = nir_tex_src_coord;
2634 p++;
2635 break;
2636 }
2637
2638 default:
2639 coord = NULL;
2640 coord_components = 0;
2641 break;
2642 }
2643
2644 switch (opcode) {
2645 case SpvOpImageSampleProjImplicitLod:
2646 case SpvOpImageSampleProjExplicitLod:
2647 case SpvOpImageSampleProjDrefImplicitLod:
2648 case SpvOpImageSampleProjDrefExplicitLod:
2649 /* These have the projector as the last coordinate component */
2650 p->src = nir_src_for_ssa(nir_channel(&b->nb, coord, coord_components));
2651 p->src_type = nir_tex_src_projector;
2652 p++;
2653 break;
2654
2655 default:
2656 break;
2657 }
2658
2659 bool is_shadow = false;
2660 unsigned gather_component = 0;
2661 switch (opcode) {
2662 case SpvOpImageSampleDrefImplicitLod:
2663 case SpvOpImageSampleDrefExplicitLod:
2664 case SpvOpImageSampleProjDrefImplicitLod:
2665 case SpvOpImageSampleProjDrefExplicitLod:
2666 case SpvOpImageDrefGather:
2667 /* These all have an explicit depth value as their next source */
2668 is_shadow = true;
2669 (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_comparator);
2670 break;
2671
2672 case SpvOpImageGather:
2673 /* This has a component as its next source */
2674 gather_component = vtn_constant_uint(b, w[idx++]);
2675 break;
2676
2677 default:
2678 break;
2679 }
2680
2681 /* For OpImageQuerySizeLod, we always have an LOD */
2682 if (opcode == SpvOpImageQuerySizeLod)
2683 (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_lod);
2684
2685 /* For OpFragmentFetchAMD, we always have a multisample index */
2686 if (opcode == SpvOpFragmentFetchAMD)
2687 (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_ms_index);
2688
2689 /* Now we need to handle some number of optional arguments */
2690 struct vtn_value *gather_offsets = NULL;
2691 if (idx < count) {
2692 uint32_t operands = w[idx];
2693
2694 if (operands & SpvImageOperandsBiasMask) {
2695 vtn_assert(texop == nir_texop_tex ||
2696 texop == nir_texop_tg4);
2697 if (texop == nir_texop_tex)
2698 texop = nir_texop_txb;
2699 uint32_t arg = image_operand_arg(b, w, count, idx,
2700 SpvImageOperandsBiasMask);
2701 (*p++) = vtn_tex_src(b, w[arg], nir_tex_src_bias);
2702 }
2703
2704 if (operands & SpvImageOperandsLodMask) {
2705 vtn_assert(texop == nir_texop_txl || texop == nir_texop_txf ||
2706 texop == nir_texop_txs || texop == nir_texop_tg4);
2707 uint32_t arg = image_operand_arg(b, w, count, idx,
2708 SpvImageOperandsLodMask);
2709 (*p++) = vtn_tex_src(b, w[arg], nir_tex_src_lod);
2710 }
2711
2712 if (operands & SpvImageOperandsGradMask) {
2713 vtn_assert(texop == nir_texop_txl);
2714 texop = nir_texop_txd;
2715 uint32_t arg = image_operand_arg(b, w, count, idx,
2716 SpvImageOperandsGradMask);
2717 (*p++) = vtn_tex_src(b, w[arg], nir_tex_src_ddx);
2718 (*p++) = vtn_tex_src(b, w[arg + 1], nir_tex_src_ddy);
2719 }
2720
2721 vtn_fail_if(util_bitcount(operands & (SpvImageOperandsConstOffsetsMask |
2722 SpvImageOperandsOffsetMask |
2723 SpvImageOperandsConstOffsetMask)) > 1,
2724 "At most one of the ConstOffset, Offset, and ConstOffsets "
2725 "image operands can be used on a given instruction.");
2726
2727 if (operands & SpvImageOperandsOffsetMask) {
2728 uint32_t arg = image_operand_arg(b, w, count, idx,
2729 SpvImageOperandsOffsetMask);
2730 (*p++) = vtn_tex_src(b, w[arg], nir_tex_src_offset);
2731 }
2732
2733 if (operands & SpvImageOperandsConstOffsetMask) {
2734 uint32_t arg = image_operand_arg(b, w, count, idx,
2735 SpvImageOperandsConstOffsetMask);
2736 (*p++) = vtn_tex_src(b, w[arg], nir_tex_src_offset);
2737 }
2738
2739 if (operands & SpvImageOperandsConstOffsetsMask) {
2740 vtn_assert(texop == nir_texop_tg4);
2741 uint32_t arg = image_operand_arg(b, w, count, idx,
2742 SpvImageOperandsConstOffsetsMask);
2743 gather_offsets = vtn_value(b, w[arg], vtn_value_type_constant);
2744 }
2745
2746 if (operands & SpvImageOperandsSampleMask) {
2747 vtn_assert(texop == nir_texop_txf_ms);
2748 uint32_t arg = image_operand_arg(b, w, count, idx,
2749 SpvImageOperandsSampleMask);
2750 texop = nir_texop_txf_ms;
2751 (*p++) = vtn_tex_src(b, w[arg], nir_tex_src_ms_index);
2752 }
2753
2754 if (operands & SpvImageOperandsMinLodMask) {
2755 vtn_assert(texop == nir_texop_tex ||
2756 texop == nir_texop_txb ||
2757 texop == nir_texop_txd);
2758 uint32_t arg = image_operand_arg(b, w, count, idx,
2759 SpvImageOperandsMinLodMask);
2760 (*p++) = vtn_tex_src(b, w[arg], nir_tex_src_min_lod);
2761 }
2762 }
2763
2764 nir_tex_instr *instr = nir_tex_instr_create(b->shader, p - srcs);
2765 instr->op = texop;
2766
2767 memcpy(instr->src, srcs, instr->num_srcs * sizeof(*instr->src));
2768
2769 instr->coord_components = coord_components;
2770 instr->sampler_dim = sampler_dim;
2771 instr->is_array = is_array;
2772 instr->is_shadow = is_shadow;
2773 instr->is_new_style_shadow =
2774 is_shadow && glsl_get_components(ret_type->type) == 1;
2775 instr->component = gather_component;
2776
2777 /* The Vulkan spec says:
2778 *
2779 * "If an instruction loads from or stores to a resource (including
2780 * atomics and image instructions) and the resource descriptor being
2781 * accessed is not dynamically uniform, then the operand corresponding
2782 * to that resource (e.g. the pointer or sampled image operand) must be
2783 * decorated with NonUniform."
2784 *
2785 * It's very careful to specify that the exact operand must be decorated
2786 * NonUniform. The SPIR-V parser is not expected to chase through long
2787 * chains to find the NonUniform decoration. It's either right there or we
2788 * can assume it doesn't exist.
2789 */
2790 enum gl_access_qualifier access = 0;
2791 vtn_foreach_decoration(b, sampled_val, non_uniform_decoration_cb, &access);
2792
2793 if (image && (access & ACCESS_NON_UNIFORM))
2794 instr->texture_non_uniform = true;
2795
2796 if (sampler && (access & ACCESS_NON_UNIFORM))
2797 instr->sampler_non_uniform = true;
2798
2799 /* for non-query ops, get dest_type from SPIR-V return type */
2800 if (dest_type == nir_type_invalid) {
2801 /* the return type should match the image type, unless the image type is
2802 * VOID (CL image), in which case the return type dictates the sampler
2803 */
2804 enum glsl_base_type sampler_base =
2805 glsl_get_sampler_result_type(image->type);
2806 enum glsl_base_type ret_base = glsl_get_base_type(ret_type->type);
2807 vtn_fail_if(sampler_base != ret_base && sampler_base != GLSL_TYPE_VOID,
2808 "SPIR-V return type mismatches image type. This is only valid "
2809 "for untyped images (OpenCL).");
2810 switch (ret_base) {
2811 case GLSL_TYPE_FLOAT: dest_type = nir_type_float; break;
2812 case GLSL_TYPE_INT: dest_type = nir_type_int; break;
2813 case GLSL_TYPE_UINT: dest_type = nir_type_uint; break;
2814 case GLSL_TYPE_BOOL: dest_type = nir_type_bool; break;
2815 default:
2816 vtn_fail("Invalid base type for sampler result");
2817 }
2818 }
2819
2820 instr->dest_type = dest_type;
2821
2822 nir_ssa_dest_init(&instr->instr, &instr->dest,
2823 nir_tex_instr_dest_size(instr), 32, NULL);
2824
2825 vtn_assert(glsl_get_vector_elements(ret_type->type) ==
2826 nir_tex_instr_dest_size(instr));
2827
2828 if (gather_offsets) {
2829 vtn_fail_if(gather_offsets->type->base_type != vtn_base_type_array ||
2830 gather_offsets->type->length != 4,
2831 "ConstOffsets must be an array of size four of vectors "
2832 "of two integer components");
2833
2834 struct vtn_type *vec_type = gather_offsets->type->array_element;
2835 vtn_fail_if(vec_type->base_type != vtn_base_type_vector ||
2836 vec_type->length != 2 ||
2837 !glsl_type_is_integer(vec_type->type),
2838 "ConstOffsets must be an array of size four of vectors "
2839 "of two integer components");
2840
2841 unsigned bit_size = glsl_get_bit_size(vec_type->type);
2842 for (uint32_t i = 0; i < 4; i++) {
2843 const nir_const_value *cvec =
2844 gather_offsets->constant->elements[i]->values;
2845 for (uint32_t j = 0; j < 2; j++) {
2846 switch (bit_size) {
2847 case 8: instr->tg4_offsets[i][j] = cvec[j].i8; break;
2848 case 16: instr->tg4_offsets[i][j] = cvec[j].i16; break;
2849 case 32: instr->tg4_offsets[i][j] = cvec[j].i32; break;
2850 case 64: instr->tg4_offsets[i][j] = cvec[j].i64; break;
2851 default:
2852 vtn_fail("Unsupported bit size: %u", bit_size);
2853 }
2854 }
2855 }
2856 }
2857
2858 nir_builder_instr_insert(&b->nb, &instr->instr);
2859
2860 vtn_push_nir_ssa(b, w[2], &instr->dest.ssa);
2861 }
2862
2863 static void
2864 fill_common_atomic_sources(struct vtn_builder *b, SpvOp opcode,
2865 const uint32_t *w, nir_src *src)
2866 {
2867 switch (opcode) {
2868 case SpvOpAtomicIIncrement:
2869 src[0] = nir_src_for_ssa(nir_imm_int(&b->nb, 1));
2870 break;
2871
2872 case SpvOpAtomicIDecrement:
2873 src[0] = nir_src_for_ssa(nir_imm_int(&b->nb, -1));
2874 break;
2875
2876 case SpvOpAtomicISub:
2877 src[0] =
2878 nir_src_for_ssa(nir_ineg(&b->nb, vtn_get_nir_ssa(b, w[6])));
2879 break;
2880
2881 case SpvOpAtomicCompareExchange:
2882 case SpvOpAtomicCompareExchangeWeak:
2883 src[0] = nir_src_for_ssa(vtn_get_nir_ssa(b, w[8]));
2884 src[1] = nir_src_for_ssa(vtn_get_nir_ssa(b, w[7]));
2885 break;
2886
2887 case SpvOpAtomicExchange:
2888 case SpvOpAtomicIAdd:
2889 case SpvOpAtomicSMin:
2890 case SpvOpAtomicUMin:
2891 case SpvOpAtomicSMax:
2892 case SpvOpAtomicUMax:
2893 case SpvOpAtomicAnd:
2894 case SpvOpAtomicOr:
2895 case SpvOpAtomicXor:
2896 case SpvOpAtomicFAddEXT:
2897 src[0] = nir_src_for_ssa(vtn_get_nir_ssa(b, w[6]));
2898 break;
2899
2900 default:
2901 vtn_fail_with_opcode("Invalid SPIR-V atomic", opcode);
2902 }
2903 }
2904
2905 static nir_ssa_def *
2906 get_image_coord(struct vtn_builder *b, uint32_t value)
2907 {
2908 nir_ssa_def *coord = vtn_get_nir_ssa(b, value);
2909
2910 /* The image_load_store intrinsics assume a 4-dim coordinate */
2911 unsigned swizzle[4];
2912 for (unsigned i = 0; i < 4; i++)
2913 swizzle[i] = MIN2(i, coord->num_components - 1);
2914
2915 return nir_swizzle(&b->nb, coord, swizzle, 4);
2916 }
2917
2918 static nir_ssa_def *
2919 expand_to_vec4(nir_builder *b, nir_ssa_def *value)
2920 {
2921 if (value->num_components == 4)
2922 return value;
2923
2924 unsigned swiz[4];
2925 for (unsigned i = 0; i < 4; i++)
2926 swiz[i] = i < value->num_components ? i : 0;
2927 return nir_swizzle(b, value, swiz, 4);
2928 }
2929
2930 static void
2931 vtn_handle_image(struct vtn_builder *b, SpvOp opcode,
2932 const uint32_t *w, unsigned count)
2933 {
2934 /* Just get this one out of the way */
2935 if (opcode == SpvOpImageTexelPointer) {
2936 struct vtn_value *val =
2937 vtn_push_value(b, w[2], vtn_value_type_image_pointer);
2938 val->image = ralloc(b, struct vtn_image_pointer);
2939
2940 val->image->image = vtn_nir_deref(b, w[3]);
2941 val->image->coord = get_image_coord(b, w[4]);
2942 val->image->sample = vtn_get_nir_ssa(b, w[5]);
2943 val->image->lod = nir_imm_int(&b->nb, 0);
2944 return;
2945 }
2946
2947 struct vtn_image_pointer image;
2948 SpvScope scope = SpvScopeInvocation;
2949 SpvMemorySemanticsMask semantics = 0;
2950
2951 enum gl_access_qualifier access = 0;
2952
2953 struct vtn_value *res_val;
2954 switch (opcode) {
2955 case SpvOpAtomicExchange:
2956 case SpvOpAtomicCompareExchange:
2957 case SpvOpAtomicCompareExchangeWeak:
2958 case SpvOpAtomicIIncrement:
2959 case SpvOpAtomicIDecrement:
2960 case SpvOpAtomicIAdd:
2961 case SpvOpAtomicISub:
2962 case SpvOpAtomicLoad:
2963 case SpvOpAtomicSMin:
2964 case SpvOpAtomicUMin:
2965 case SpvOpAtomicSMax:
2966 case SpvOpAtomicUMax:
2967 case SpvOpAtomicAnd:
2968 case SpvOpAtomicOr:
2969 case SpvOpAtomicXor:
2970 case SpvOpAtomicFAddEXT:
2971 res_val = vtn_value(b, w[3], vtn_value_type_image_pointer);
2972 image = *res_val->image;
2973 scope = vtn_constant_uint(b, w[4]);
2974 semantics = vtn_constant_uint(b, w[5]);
2975 access |= ACCESS_COHERENT;
2976 break;
2977
2978 case SpvOpAtomicStore:
2979 res_val = vtn_value(b, w[1], vtn_value_type_image_pointer);
2980 image = *res_val->image;
2981 scope = vtn_constant_uint(b, w[2]);
2982 semantics = vtn_constant_uint(b, w[3]);
2983 access |= ACCESS_COHERENT;
2984 break;
2985
2986 case SpvOpImageQuerySizeLod:
2987 res_val = vtn_untyped_value(b, w[3]);
2988 image.image = vtn_get_image(b, w[3]);
2989 image.coord = NULL;
2990 image.sample = NULL;
2991 image.lod = vtn_ssa_value(b, w[4])->def;
2992 break;
2993
2994 case SpvOpImageQuerySize:
2995 res_val = vtn_untyped_value(b, w[3]);
2996 image.image = vtn_get_image(b, w[3]);
2997 image.coord = NULL;
2998 image.sample = NULL;
2999 image.lod = NULL;
3000 break;
3001
3002 case SpvOpImageQueryFormat:
3003 case SpvOpImageQueryOrder:
3004 res_val = vtn_untyped_value(b, w[3]);
3005 image.image = vtn_get_image(b, w[3]);
3006 image.coord = NULL;
3007 image.sample = NULL;
3008 image.lod = NULL;
3009 break;
3010
3011 case SpvOpImageRead: {
3012 res_val = vtn_untyped_value(b, w[3]);
3013 image.image = vtn_get_image(b, w[3]);
3014 image.coord = get_image_coord(b, w[4]);
3015
3016 const SpvImageOperandsMask operands =
3017 count > 5 ? w[5] : SpvImageOperandsMaskNone;
3018
3019 if (operands & SpvImageOperandsSampleMask) {
3020 uint32_t arg = image_operand_arg(b, w, count, 5,
3021 SpvImageOperandsSampleMask);
3022 image.sample = vtn_get_nir_ssa(b, w[arg]);
3023 } else {
3024 image.sample = nir_ssa_undef(&b->nb, 1, 32);
3025 }
3026
3027 if (operands & SpvImageOperandsMakeTexelVisibleMask) {
3028 vtn_fail_if((operands & SpvImageOperandsNonPrivateTexelMask) == 0,
3029 "MakeTexelVisible requires NonPrivateTexel to also be set.");
3030 uint32_t arg = image_operand_arg(b, w, count, 5,
3031 SpvImageOperandsMakeTexelVisibleMask);
3032 semantics = SpvMemorySemanticsMakeVisibleMask;
3033 scope = vtn_constant_uint(b, w[arg]);
3034 }
3035
3036 if (operands & SpvImageOperandsLodMask) {
3037 uint32_t arg = image_operand_arg(b, w, count, 5,
3038 SpvImageOperandsLodMask);
3039 image.lod = vtn_get_nir_ssa(b, w[arg]);
3040 } else {
3041 image.lod = nir_imm_int(&b->nb, 0);
3042 }
3043
3044 if (operands & SpvImageOperandsVolatileTexelMask)
3045 access |= ACCESS_VOLATILE;
3046
3047 break;
3048 }
3049
3050 case SpvOpImageWrite: {
3051 res_val = vtn_untyped_value(b, w[1]);
3052 image.image = vtn_get_image(b, w[1]);
3053 image.coord = get_image_coord(b, w[2]);
3054
3055 /* texel = w[3] */
3056
3057 const SpvImageOperandsMask operands =
3058 count > 4 ? w[4] : SpvImageOperandsMaskNone;
3059
3060 if (operands & SpvImageOperandsSampleMask) {
3061 uint32_t arg = image_operand_arg(b, w, count, 4,
3062 SpvImageOperandsSampleMask);
3063 image.sample = vtn_get_nir_ssa(b, w[arg]);
3064 } else {
3065 image.sample = nir_ssa_undef(&b->nb, 1, 32);
3066 }
3067
3068 if (operands & SpvImageOperandsMakeTexelAvailableMask) {
3069 vtn_fail_if((operands & SpvImageOperandsNonPrivateTexelMask) == 0,
3070 "MakeTexelAvailable requires NonPrivateTexel to also be set.");
3071 uint32_t arg = image_operand_arg(b, w, count, 4,
3072 SpvImageOperandsMakeTexelAvailableMask);
3073 semantics = SpvMemorySemanticsMakeAvailableMask;
3074 scope = vtn_constant_uint(b, w[arg]);
3075 }
3076
3077 if (operands & SpvImageOperandsLodMask) {
3078 uint32_t arg = image_operand_arg(b, w, count, 4,
3079 SpvImageOperandsLodMask);
3080 image.lod = vtn_get_nir_ssa(b, w[arg]);
3081 } else {
3082 image.lod = nir_imm_int(&b->nb, 0);
3083 }
3084
3085 if (operands & SpvImageOperandsVolatileTexelMask)
3086 access |= ACCESS_VOLATILE;
3087
3088 break;
3089 }
3090
3091 default:
3092 vtn_fail_with_opcode("Invalid image opcode", opcode);
3093 }
3094
3095 if (semantics & SpvMemorySemanticsVolatileMask)
3096 access |= ACCESS_VOLATILE;
3097
3098 nir_intrinsic_op op;
3099 switch (opcode) {
3100 #define OP(S, N) case SpvOp##S: op = nir_intrinsic_image_deref_##N; break;
3101 OP(ImageQuerySize, size)
3102 OP(ImageQuerySizeLod, size)
3103 OP(ImageRead, load)
3104 OP(ImageWrite, store)
3105 OP(AtomicLoad, load)
3106 OP(AtomicStore, store)
3107 OP(AtomicExchange, atomic_exchange)
3108 OP(AtomicCompareExchange, atomic_comp_swap)
3109 OP(AtomicCompareExchangeWeak, atomic_comp_swap)
3110 OP(AtomicIIncrement, atomic_add)
3111 OP(AtomicIDecrement, atomic_add)
3112 OP(AtomicIAdd, atomic_add)
3113 OP(AtomicISub, atomic_add)
3114 OP(AtomicSMin, atomic_imin)
3115 OP(AtomicUMin, atomic_umin)
3116 OP(AtomicSMax, atomic_imax)
3117 OP(AtomicUMax, atomic_umax)
3118 OP(AtomicAnd, atomic_and)
3119 OP(AtomicOr, atomic_or)
3120 OP(AtomicXor, atomic_xor)
3121 OP(AtomicFAddEXT, atomic_fadd)
3122 OP(ImageQueryFormat, format)
3123 OP(ImageQueryOrder, order)
3124 #undef OP
3125 default:
3126 vtn_fail_with_opcode("Invalid image opcode", opcode);
3127 }
3128
3129 nir_intrinsic_instr *intrin = nir_intrinsic_instr_create(b->shader, op);
3130
3131 intrin->src[0] = nir_src_for_ssa(&image.image->dest.ssa);
3132
3133 switch (opcode) {
3134 case SpvOpImageQuerySize:
3135 case SpvOpImageQuerySizeLod:
3136 case SpvOpImageQueryFormat:
3137 case SpvOpImageQueryOrder:
3138 break;
3139 default:
3140 /* The image coordinate is always 4 components but we may not have that
3141 * many. Swizzle to compensate.
3142 */
3143 intrin->src[1] = nir_src_for_ssa(expand_to_vec4(&b->nb, image.coord));
3144 intrin->src[2] = nir_src_for_ssa(image.sample);
3145 break;
3146 }
3147
3148 /* The Vulkan spec says:
3149 *
3150 * "If an instruction loads from or stores to a resource (including
3151 * atomics and image instructions) and the resource descriptor being
3152 * accessed is not dynamically uniform, then the operand corresponding
3153 * to that resource (e.g. the pointer or sampled image operand) must be
3154 * decorated with NonUniform."
3155 *
3156 * It's very careful to specify that the exact operand must be decorated
3157 * NonUniform. The SPIR-V parser is not expected to chase through long
3158 * chains to find the NonUniform decoration. It's either right there or we
3159 * can assume it doesn't exist.
3160 */
3161 vtn_foreach_decoration(b, res_val, non_uniform_decoration_cb, &access);
3162 nir_intrinsic_set_access(intrin, access);
3163
3164 switch (opcode) {
3165 case SpvOpImageQueryFormat:
3166 case SpvOpImageQueryOrder:
3167 /* No additional sources */
3168 break;
3169 case SpvOpImageQuerySize:
3170 intrin->src[1] = nir_src_for_ssa(nir_imm_int(&b->nb, 0));
3171 break;
3172 case SpvOpImageQuerySizeLod:
3173 intrin->src[1] = nir_src_for_ssa(image.lod);
3174 break;
3175 case SpvOpAtomicLoad:
3176 case SpvOpImageRead:
3177 /* Only OpImageRead can support a lod parameter if
3178 * SPV_AMD_shader_image_load_store_lod is used but the current NIR
3179 * intrinsics definition for atomics requires us to set it for
3180 * OpAtomicLoad.
3181 */
3182 intrin->src[3] = nir_src_for_ssa(image.lod);
3183 break;
3184 case SpvOpAtomicStore:
3185 case SpvOpImageWrite: {
3186 const uint32_t value_id = opcode == SpvOpAtomicStore ? w[4] : w[3];
3187 struct vtn_ssa_value *value = vtn_ssa_value(b, value_id);
3188 /* nir_intrinsic_image_deref_store always takes a vec4 value */
3189 assert(op == nir_intrinsic_image_deref_store);
3190 intrin->num_components = 4;
3191 intrin->src[3] = nir_src_for_ssa(expand_to_vec4(&b->nb, value->def));
3192 /* Only OpImageWrite can support a lod parameter if
3193 * SPV_AMD_shader_image_load_store_lod is used but the current NIR
3194 * intrinsics definition for atomics requires us to set it for
3195 * OpAtomicStore.
3196 */
3197 intrin->src[4] = nir_src_for_ssa(image.lod);
3198
3199 if (opcode == SpvOpImageWrite)
3200 nir_intrinsic_set_type(intrin, nir_get_nir_type_for_glsl_type(value->type));
3201 break;
3202 }
3203
3204 case SpvOpAtomicCompareExchange:
3205 case SpvOpAtomicCompareExchangeWeak:
3206 case SpvOpAtomicIIncrement:
3207 case SpvOpAtomicIDecrement:
3208 case SpvOpAtomicExchange:
3209 case SpvOpAtomicIAdd:
3210 case SpvOpAtomicISub:
3211 case SpvOpAtomicSMin:
3212 case SpvOpAtomicUMin:
3213 case SpvOpAtomicSMax:
3214 case SpvOpAtomicUMax:
3215 case SpvOpAtomicAnd:
3216 case SpvOpAtomicOr:
3217 case SpvOpAtomicXor:
3218 case SpvOpAtomicFAddEXT:
3219 fill_common_atomic_sources(b, opcode, w, &intrin->src[3]);
3220 break;
3221
3222 default:
3223 vtn_fail_with_opcode("Invalid image opcode", opcode);
3224 }
3225
3226 /* Image operations implicitly have the Image storage memory semantics. */
3227 semantics |= SpvMemorySemanticsImageMemoryMask;
3228
3229 SpvMemorySemanticsMask before_semantics;
3230 SpvMemorySemanticsMask after_semantics;
3231 vtn_split_barrier_semantics(b, semantics, &before_semantics, &after_semantics);
3232
3233 if (before_semantics)
3234 vtn_emit_memory_barrier(b, scope, before_semantics);
3235
3236 if (opcode != SpvOpImageWrite && opcode != SpvOpAtomicStore) {
3237 struct vtn_type *type = vtn_get_type(b, w[1]);
3238
3239 unsigned dest_components = glsl_get_vector_elements(type->type);
3240 if (nir_intrinsic_infos[op].dest_components == 0)
3241 intrin->num_components = dest_components;
3242
3243 nir_ssa_dest_init(&intrin->instr, &intrin->dest,
3244 nir_intrinsic_dest_components(intrin), 32, NULL);
3245
3246 nir_builder_instr_insert(&b->nb, &intrin->instr);
3247
3248 nir_ssa_def *result = &intrin->dest.ssa;
3249 if (nir_intrinsic_dest_components(intrin) != dest_components)
3250 result = nir_channels(&b->nb, result, (1 << dest_components) - 1);
3251
3252 vtn_push_nir_ssa(b, w[2], result);
3253
3254 if (opcode == SpvOpImageRead)
3255 nir_intrinsic_set_type(intrin, nir_get_nir_type_for_glsl_type(type->type));
3256 } else {
3257 nir_builder_instr_insert(&b->nb, &intrin->instr);
3258 }
3259
3260 if (after_semantics)
3261 vtn_emit_memory_barrier(b, scope, after_semantics);
3262 }
3263
3264 static nir_intrinsic_op
3265 get_ssbo_nir_atomic_op(struct vtn_builder *b, SpvOp opcode)
3266 {
3267 switch (opcode) {
3268 case SpvOpAtomicLoad: return nir_intrinsic_load_ssbo;
3269 case SpvOpAtomicStore: return nir_intrinsic_store_ssbo;
3270 #define OP(S, N) case SpvOp##S: return nir_intrinsic_ssbo_##N;
3271 OP(AtomicExchange, atomic_exchange)
3272 OP(AtomicCompareExchange, atomic_comp_swap)
3273 OP(AtomicCompareExchangeWeak, atomic_comp_swap)
3274 OP(AtomicIIncrement, atomic_add)
3275 OP(AtomicIDecrement, atomic_add)
3276 OP(AtomicIAdd, atomic_add)
3277 OP(AtomicISub, atomic_add)
3278 OP(AtomicSMin, atomic_imin)
3279 OP(AtomicUMin, atomic_umin)
3280 OP(AtomicSMax, atomic_imax)
3281 OP(AtomicUMax, atomic_umax)
3282 OP(AtomicAnd, atomic_and)
3283 OP(AtomicOr, atomic_or)
3284 OP(AtomicXor, atomic_xor)
3285 OP(AtomicFAddEXT, atomic_fadd)
3286 #undef OP
3287 default:
3288 vtn_fail_with_opcode("Invalid SSBO atomic", opcode);
3289 }
3290 }
3291
3292 static nir_intrinsic_op
3293 get_uniform_nir_atomic_op(struct vtn_builder *b, SpvOp opcode)
3294 {
3295 switch (opcode) {
3296 #define OP(S, N) case SpvOp##S: return nir_intrinsic_atomic_counter_ ##N;
3297 OP(AtomicLoad, read_deref)
3298 OP(AtomicExchange, exchange)
3299 OP(AtomicCompareExchange, comp_swap)
3300 OP(AtomicCompareExchangeWeak, comp_swap)
3301 OP(AtomicIIncrement, inc_deref)
3302 OP(AtomicIDecrement, post_dec_deref)
3303 OP(AtomicIAdd, add_deref)
3304 OP(AtomicISub, add_deref)
3305 OP(AtomicUMin, min_deref)
3306 OP(AtomicUMax, max_deref)
3307 OP(AtomicAnd, and_deref)
3308 OP(AtomicOr, or_deref)
3309 OP(AtomicXor, xor_deref)
3310 #undef OP
3311 default:
3312 /* We left the following out: AtomicStore, AtomicSMin and
3313 * AtomicSmax. Right now there are not nir intrinsics for them. At this
3314 * moment Atomic Counter support is needed for ARB_spirv support, so is
3315 * only need to support GLSL Atomic Counters that are uints and don't
3316 * allow direct storage.
3317 */
3318 vtn_fail("Invalid uniform atomic");
3319 }
3320 }
3321
3322 static nir_intrinsic_op
3323 get_deref_nir_atomic_op(struct vtn_builder *b, SpvOp opcode)
3324 {
3325 switch (opcode) {
3326 case SpvOpAtomicLoad: return nir_intrinsic_load_deref;
3327 case SpvOpAtomicStore: return nir_intrinsic_store_deref;
3328 #define OP(S, N) case SpvOp##S: return nir_intrinsic_deref_##N;
3329 OP(AtomicExchange, atomic_exchange)
3330 OP(AtomicCompareExchange, atomic_comp_swap)
3331 OP(AtomicCompareExchangeWeak, atomic_comp_swap)
3332 OP(AtomicIIncrement, atomic_add)
3333 OP(AtomicIDecrement, atomic_add)
3334 OP(AtomicIAdd, atomic_add)
3335 OP(AtomicISub, atomic_add)
3336 OP(AtomicSMin, atomic_imin)
3337 OP(AtomicUMin, atomic_umin)
3338 OP(AtomicSMax, atomic_imax)
3339 OP(AtomicUMax, atomic_umax)
3340 OP(AtomicAnd, atomic_and)
3341 OP(AtomicOr, atomic_or)
3342 OP(AtomicXor, atomic_xor)
3343 OP(AtomicFAddEXT, atomic_fadd)
3344 #undef OP
3345 default:
3346 vtn_fail_with_opcode("Invalid shared atomic", opcode);
3347 }
3348 }
3349
3350 /*
3351 * Handles shared atomics, ssbo atomics and atomic counters.
3352 */
3353 static void
3354 vtn_handle_atomics(struct vtn_builder *b, SpvOp opcode,
3355 const uint32_t *w, UNUSED unsigned count)
3356 {
3357 struct vtn_pointer *ptr;
3358 nir_intrinsic_instr *atomic;
3359
3360 SpvScope scope = SpvScopeInvocation;
3361 SpvMemorySemanticsMask semantics = 0;
3362 enum gl_access_qualifier access = 0;
3363
3364 switch (opcode) {
3365 case SpvOpAtomicLoad:
3366 case SpvOpAtomicExchange:
3367 case SpvOpAtomicCompareExchange:
3368 case SpvOpAtomicCompareExchangeWeak:
3369 case SpvOpAtomicIIncrement:
3370 case SpvOpAtomicIDecrement:
3371 case SpvOpAtomicIAdd:
3372 case SpvOpAtomicISub:
3373 case SpvOpAtomicSMin:
3374 case SpvOpAtomicUMin:
3375 case SpvOpAtomicSMax:
3376 case SpvOpAtomicUMax:
3377 case SpvOpAtomicAnd:
3378 case SpvOpAtomicOr:
3379 case SpvOpAtomicXor:
3380 case SpvOpAtomicFAddEXT:
3381 ptr = vtn_value(b, w[3], vtn_value_type_pointer)->pointer;
3382 scope = vtn_constant_uint(b, w[4]);
3383 semantics = vtn_constant_uint(b, w[5]);
3384 break;
3385
3386 case SpvOpAtomicStore:
3387 ptr = vtn_value(b, w[1], vtn_value_type_pointer)->pointer;
3388 scope = vtn_constant_uint(b, w[2]);
3389 semantics = vtn_constant_uint(b, w[3]);
3390 break;
3391
3392 default:
3393 vtn_fail_with_opcode("Invalid SPIR-V atomic", opcode);
3394 }
3395
3396 if (semantics & SpvMemorySemanticsVolatileMask)
3397 access |= ACCESS_VOLATILE;
3398
3399 /* uniform as "atomic counter uniform" */
3400 if (ptr->mode == vtn_variable_mode_atomic_counter) {
3401 nir_deref_instr *deref = vtn_pointer_to_deref(b, ptr);
3402 nir_intrinsic_op op = get_uniform_nir_atomic_op(b, opcode);
3403 atomic = nir_intrinsic_instr_create(b->nb.shader, op);
3404 atomic->src[0] = nir_src_for_ssa(&deref->dest.ssa);
3405
3406 /* SSBO needs to initialize index/offset. In this case we don't need to,
3407 * as that info is already stored on the ptr->var->var nir_variable (see
3408 * vtn_create_variable)
3409 */
3410
3411 switch (opcode) {
3412 case SpvOpAtomicLoad:
3413 case SpvOpAtomicExchange:
3414 case SpvOpAtomicCompareExchange:
3415 case SpvOpAtomicCompareExchangeWeak:
3416 case SpvOpAtomicIIncrement:
3417 case SpvOpAtomicIDecrement:
3418 case SpvOpAtomicIAdd:
3419 case SpvOpAtomicISub:
3420 case SpvOpAtomicSMin:
3421 case SpvOpAtomicUMin:
3422 case SpvOpAtomicSMax:
3423 case SpvOpAtomicUMax:
3424 case SpvOpAtomicAnd:
3425 case SpvOpAtomicOr:
3426 case SpvOpAtomicXor:
3427 /* Nothing: we don't need to call fill_common_atomic_sources here, as
3428 * atomic counter uniforms doesn't have sources
3429 */
3430 break;
3431
3432 default:
3433 unreachable("Invalid SPIR-V atomic");
3434
3435 }
3436 } else if (vtn_pointer_uses_ssa_offset(b, ptr)) {
3437 nir_ssa_def *offset, *index;
3438 offset = vtn_pointer_to_offset(b, ptr, &index);
3439
3440 assert(ptr->mode == vtn_variable_mode_ssbo);
3441
3442 nir_intrinsic_op op = get_ssbo_nir_atomic_op(b, opcode);
3443 atomic = nir_intrinsic_instr_create(b->nb.shader, op);
3444
3445 nir_intrinsic_set_access(atomic, access | ACCESS_COHERENT);
3446
3447 int src = 0;
3448 switch (opcode) {
3449 case SpvOpAtomicLoad:
3450 atomic->num_components = glsl_get_vector_elements(ptr->type->type);
3451 nir_intrinsic_set_align(atomic, 4, 0);
3452 if (ptr->mode == vtn_variable_mode_ssbo)
3453 atomic->src[src++] = nir_src_for_ssa(index);
3454 atomic->src[src++] = nir_src_for_ssa(offset);
3455 break;
3456
3457 case SpvOpAtomicStore:
3458 atomic->num_components = glsl_get_vector_elements(ptr->type->type);
3459 nir_intrinsic_set_write_mask(atomic, (1 << atomic->num_components) - 1);
3460 nir_intrinsic_set_align(atomic, 4, 0);
3461 atomic->src[src++] = nir_src_for_ssa(vtn_get_nir_ssa(b, w[4]));
3462 if (ptr->mode == vtn_variable_mode_ssbo)
3463 atomic->src[src++] = nir_src_for_ssa(index);
3464 atomic->src[src++] = nir_src_for_ssa(offset);
3465 break;
3466
3467 case SpvOpAtomicExchange:
3468 case SpvOpAtomicCompareExchange:
3469 case SpvOpAtomicCompareExchangeWeak:
3470 case SpvOpAtomicIIncrement:
3471 case SpvOpAtomicIDecrement:
3472 case SpvOpAtomicIAdd:
3473 case SpvOpAtomicISub:
3474 case SpvOpAtomicSMin:
3475 case SpvOpAtomicUMin:
3476 case SpvOpAtomicSMax:
3477 case SpvOpAtomicUMax:
3478 case SpvOpAtomicAnd:
3479 case SpvOpAtomicOr:
3480 case SpvOpAtomicXor:
3481 case SpvOpAtomicFAddEXT:
3482 if (ptr->mode == vtn_variable_mode_ssbo)
3483 atomic->src[src++] = nir_src_for_ssa(index);
3484 atomic->src[src++] = nir_src_for_ssa(offset);
3485 fill_common_atomic_sources(b, opcode, w, &atomic->src[src]);
3486 break;
3487
3488 default:
3489 vtn_fail_with_opcode("Invalid SPIR-V atomic", opcode);
3490 }
3491 } else {
3492 nir_deref_instr *deref = vtn_pointer_to_deref(b, ptr);
3493 const struct glsl_type *deref_type = deref->type;
3494 nir_intrinsic_op op = get_deref_nir_atomic_op(b, opcode);
3495 atomic = nir_intrinsic_instr_create(b->nb.shader, op);
3496 atomic->src[0] = nir_src_for_ssa(&deref->dest.ssa);
3497
3498 if (ptr->mode != vtn_variable_mode_workgroup)
3499 access |= ACCESS_COHERENT;
3500
3501 nir_intrinsic_set_access(atomic, access);
3502
3503 switch (opcode) {
3504 case SpvOpAtomicLoad:
3505 atomic->num_components = glsl_get_vector_elements(deref_type);
3506 break;
3507
3508 case SpvOpAtomicStore:
3509 atomic->num_components = glsl_get_vector_elements(deref_type);
3510 nir_intrinsic_set_write_mask(atomic, (1 << atomic->num_components) - 1);
3511 atomic->src[1] = nir_src_for_ssa(vtn_get_nir_ssa(b, w[4]));
3512 break;
3513
3514 case SpvOpAtomicExchange:
3515 case SpvOpAtomicCompareExchange:
3516 case SpvOpAtomicCompareExchangeWeak:
3517 case SpvOpAtomicIIncrement:
3518 case SpvOpAtomicIDecrement:
3519 case SpvOpAtomicIAdd:
3520 case SpvOpAtomicISub:
3521 case SpvOpAtomicSMin:
3522 case SpvOpAtomicUMin:
3523 case SpvOpAtomicSMax:
3524 case SpvOpAtomicUMax:
3525 case SpvOpAtomicAnd:
3526 case SpvOpAtomicOr:
3527 case SpvOpAtomicXor:
3528 case SpvOpAtomicFAddEXT:
3529 fill_common_atomic_sources(b, opcode, w, &atomic->src[1]);
3530 break;
3531
3532 default:
3533 vtn_fail_with_opcode("Invalid SPIR-V atomic", opcode);
3534 }
3535 }
3536
3537 /* Atomic ordering operations will implicitly apply to the atomic operation
3538 * storage class, so include that too.
3539 */
3540 semantics |= vtn_mode_to_memory_semantics(ptr->mode);
3541
3542 SpvMemorySemanticsMask before_semantics;
3543 SpvMemorySemanticsMask after_semantics;
3544 vtn_split_barrier_semantics(b, semantics, &before_semantics, &after_semantics);
3545
3546 if (before_semantics)
3547 vtn_emit_memory_barrier(b, scope, before_semantics);
3548
3549 if (opcode != SpvOpAtomicStore) {
3550 struct vtn_type *type = vtn_get_type(b, w[1]);
3551
3552 nir_ssa_dest_init(&atomic->instr, &atomic->dest,
3553 glsl_get_vector_elements(type->type),
3554 glsl_get_bit_size(type->type), NULL);
3555
3556 vtn_push_nir_ssa(b, w[2], &atomic->dest.ssa);
3557 }
3558
3559 nir_builder_instr_insert(&b->nb, &atomic->instr);
3560
3561 if (after_semantics)
3562 vtn_emit_memory_barrier(b, scope, after_semantics);
3563 }
3564
3565 static nir_alu_instr *
3566 create_vec(struct vtn_builder *b, unsigned num_components, unsigned bit_size)
3567 {
3568 nir_op op = nir_op_vec(num_components);
3569 nir_alu_instr *vec = nir_alu_instr_create(b->shader, op);
3570 nir_ssa_dest_init(&vec->instr, &vec->dest.dest, num_components,
3571 bit_size, NULL);
3572 vec->dest.write_mask = (1 <