604174bcaaa7007e3d5421ce81c4b280fcf2f026
[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 vtn_fail_if(sampled_type->base_type != vtn_base_type_scalar ||
1573 glsl_get_bit_size(sampled_type->type) != 32,
1574 "Sampled type of OpTypeImage must be a 32-bit scalar");
1575
1576 enum glsl_sampler_dim dim;
1577 switch ((SpvDim)w[3]) {
1578 case SpvDim1D: dim = GLSL_SAMPLER_DIM_1D; break;
1579 case SpvDim2D: dim = GLSL_SAMPLER_DIM_2D; break;
1580 case SpvDim3D: dim = GLSL_SAMPLER_DIM_3D; break;
1581 case SpvDimCube: dim = GLSL_SAMPLER_DIM_CUBE; break;
1582 case SpvDimRect: dim = GLSL_SAMPLER_DIM_RECT; break;
1583 case SpvDimBuffer: dim = GLSL_SAMPLER_DIM_BUF; break;
1584 case SpvDimSubpassData: dim = GLSL_SAMPLER_DIM_SUBPASS; break;
1585 default:
1586 vtn_fail("Invalid SPIR-V image dimensionality: %s (%u)",
1587 spirv_dim_to_string((SpvDim)w[3]), w[3]);
1588 }
1589
1590 /* w[4]: as per Vulkan spec "Validation Rules within a Module",
1591 * The “Depth” operand of OpTypeImage is ignored.
1592 */
1593 bool is_array = w[5];
1594 bool multisampled = w[6];
1595 unsigned sampled = w[7];
1596 SpvImageFormat format = w[8];
1597
1598 if (count > 9)
1599 val->type->access_qualifier = w[9];
1600 else
1601 val->type->access_qualifier = SpvAccessQualifierReadWrite;
1602
1603 if (multisampled) {
1604 if (dim == GLSL_SAMPLER_DIM_2D)
1605 dim = GLSL_SAMPLER_DIM_MS;
1606 else if (dim == GLSL_SAMPLER_DIM_SUBPASS)
1607 dim = GLSL_SAMPLER_DIM_SUBPASS_MS;
1608 else
1609 vtn_fail("Unsupported multisampled image type");
1610 }
1611
1612 val->type->image_format = translate_image_format(b, format);
1613
1614 enum glsl_base_type sampled_base_type =
1615 glsl_get_base_type(sampled_type->type);
1616 if (sampled == 1) {
1617 val->type->glsl_image = glsl_sampler_type(dim, false, is_array,
1618 sampled_base_type);
1619 } else if (sampled == 2) {
1620 val->type->glsl_image = glsl_image_type(dim, is_array,
1621 sampled_base_type);
1622 } else {
1623 vtn_fail("We need to know if the image will be sampled");
1624 }
1625 break;
1626 }
1627
1628 case SpvOpTypeSampledImage: {
1629 val->type->base_type = vtn_base_type_sampled_image;
1630 val->type->image = vtn_get_type(b, w[2]);
1631
1632 /* Sampled images are represented NIR as a vec2 SSA value where each
1633 * component is the result of a deref instruction. The first component
1634 * is the image and the second is the sampler. An OpLoad on an
1635 * OpTypeSampledImage pointer from UniformConstant memory just takes
1636 * the NIR deref from the pointer and duplicates it to both vector
1637 * components.
1638 */
1639 nir_address_format addr_format =
1640 vtn_mode_to_address_format(b, vtn_variable_mode_function);
1641 assert(nir_address_format_num_components(addr_format) == 1);
1642 unsigned bit_size = nir_address_format_bit_size(addr_format);
1643 assert(bit_size == 32 || bit_size == 64);
1644
1645 enum glsl_base_type base_type =
1646 bit_size == 32 ? GLSL_TYPE_UINT : GLSL_TYPE_UINT64;
1647 val->type->type = glsl_vector_type(base_type, 2);
1648 break;
1649 }
1650
1651 case SpvOpTypeSampler:
1652 val->type->base_type = vtn_base_type_sampler;
1653
1654 /* Samplers are represented in NIR as a scalar SSA value that is the
1655 * result of a deref instruction. An OpLoad on an OpTypeSampler pointer
1656 * from UniformConstant memory just takes the NIR deref from the pointer
1657 * and turns it into an SSA value.
1658 */
1659 val->type->type = nir_address_format_to_glsl_type(
1660 vtn_mode_to_address_format(b, vtn_variable_mode_function));
1661 break;
1662
1663 case SpvOpTypeOpaque:
1664 case SpvOpTypeEvent:
1665 case SpvOpTypeDeviceEvent:
1666 case SpvOpTypeReserveId:
1667 case SpvOpTypeQueue:
1668 case SpvOpTypePipe:
1669 default:
1670 vtn_fail_with_opcode("Unhandled opcode", opcode);
1671 }
1672
1673 vtn_foreach_decoration(b, val, type_decoration_cb, NULL);
1674
1675 if (val->type->base_type == vtn_base_type_struct &&
1676 (val->type->block || val->type->buffer_block)) {
1677 for (unsigned i = 0; i < val->type->length; i++) {
1678 vtn_fail_if(vtn_type_contains_block(b, val->type->members[i]),
1679 "Block and BufferBlock decorations cannot decorate a "
1680 "structure type that is nested at any level inside "
1681 "another structure type decorated with Block or "
1682 "BufferBlock.");
1683 }
1684 }
1685 }
1686
1687 static nir_constant *
1688 vtn_null_constant(struct vtn_builder *b, struct vtn_type *type)
1689 {
1690 nir_constant *c = rzalloc(b, nir_constant);
1691
1692 switch (type->base_type) {
1693 case vtn_base_type_scalar:
1694 case vtn_base_type_vector:
1695 /* Nothing to do here. It's already initialized to zero */
1696 break;
1697
1698 case vtn_base_type_pointer: {
1699 enum vtn_variable_mode mode = vtn_storage_class_to_mode(
1700 b, type->storage_class, type->deref, NULL);
1701 nir_address_format addr_format = vtn_mode_to_address_format(b, mode);
1702
1703 const nir_const_value *null_value = nir_address_format_null_value(addr_format);
1704 memcpy(c->values, null_value,
1705 sizeof(nir_const_value) * nir_address_format_num_components(addr_format));
1706 break;
1707 }
1708
1709 case vtn_base_type_void:
1710 case vtn_base_type_image:
1711 case vtn_base_type_sampler:
1712 case vtn_base_type_sampled_image:
1713 case vtn_base_type_function:
1714 /* For those we have to return something but it doesn't matter what. */
1715 break;
1716
1717 case vtn_base_type_matrix:
1718 case vtn_base_type_array:
1719 vtn_assert(type->length > 0);
1720 c->num_elements = type->length;
1721 c->elements = ralloc_array(b, nir_constant *, c->num_elements);
1722
1723 c->elements[0] = vtn_null_constant(b, type->array_element);
1724 for (unsigned i = 1; i < c->num_elements; i++)
1725 c->elements[i] = c->elements[0];
1726 break;
1727
1728 case vtn_base_type_struct:
1729 c->num_elements = type->length;
1730 c->elements = ralloc_array(b, nir_constant *, c->num_elements);
1731 for (unsigned i = 0; i < c->num_elements; i++)
1732 c->elements[i] = vtn_null_constant(b, type->members[i]);
1733 break;
1734
1735 default:
1736 vtn_fail("Invalid type for null constant");
1737 }
1738
1739 return c;
1740 }
1741
1742 static void
1743 spec_constant_decoration_cb(struct vtn_builder *b, UNUSED struct vtn_value *val,
1744 ASSERTED int member,
1745 const struct vtn_decoration *dec, void *data)
1746 {
1747 vtn_assert(member == -1);
1748 if (dec->decoration != SpvDecorationSpecId)
1749 return;
1750
1751 nir_const_value *value = data;
1752 for (unsigned i = 0; i < b->num_specializations; i++) {
1753 if (b->specializations[i].id == dec->operands[0]) {
1754 *value = b->specializations[i].value;
1755 return;
1756 }
1757 }
1758 }
1759
1760 static void
1761 handle_workgroup_size_decoration_cb(struct vtn_builder *b,
1762 struct vtn_value *val,
1763 ASSERTED int member,
1764 const struct vtn_decoration *dec,
1765 UNUSED void *data)
1766 {
1767 vtn_assert(member == -1);
1768 if (dec->decoration != SpvDecorationBuiltIn ||
1769 dec->operands[0] != SpvBuiltInWorkgroupSize)
1770 return;
1771
1772 vtn_assert(val->type->type == glsl_vector_type(GLSL_TYPE_UINT, 3));
1773 b->workgroup_size_builtin = val;
1774 }
1775
1776 static void
1777 vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
1778 const uint32_t *w, unsigned count)
1779 {
1780 struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_constant);
1781 val->constant = rzalloc(b, nir_constant);
1782 switch (opcode) {
1783 case SpvOpConstantTrue:
1784 case SpvOpConstantFalse:
1785 case SpvOpSpecConstantTrue:
1786 case SpvOpSpecConstantFalse: {
1787 vtn_fail_if(val->type->type != glsl_bool_type(),
1788 "Result type of %s must be OpTypeBool",
1789 spirv_op_to_string(opcode));
1790
1791 bool bval = (opcode == SpvOpConstantTrue ||
1792 opcode == SpvOpSpecConstantTrue);
1793
1794 nir_const_value u32val = nir_const_value_for_uint(bval, 32);
1795
1796 if (opcode == SpvOpSpecConstantTrue ||
1797 opcode == SpvOpSpecConstantFalse)
1798 vtn_foreach_decoration(b, val, spec_constant_decoration_cb, &u32val);
1799
1800 val->constant->values[0].b = u32val.u32 != 0;
1801 break;
1802 }
1803
1804 case SpvOpConstant:
1805 case SpvOpSpecConstant: {
1806 vtn_fail_if(val->type->base_type != vtn_base_type_scalar,
1807 "Result type of %s must be a scalar",
1808 spirv_op_to_string(opcode));
1809 int bit_size = glsl_get_bit_size(val->type->type);
1810 switch (bit_size) {
1811 case 64:
1812 val->constant->values[0].u64 = vtn_u64_literal(&w[3]);
1813 break;
1814 case 32:
1815 val->constant->values[0].u32 = w[3];
1816 break;
1817 case 16:
1818 val->constant->values[0].u16 = w[3];
1819 break;
1820 case 8:
1821 val->constant->values[0].u8 = w[3];
1822 break;
1823 default:
1824 vtn_fail("Unsupported SpvOpConstant bit size: %u", bit_size);
1825 }
1826
1827 if (opcode == SpvOpSpecConstant)
1828 vtn_foreach_decoration(b, val, spec_constant_decoration_cb,
1829 &val->constant->values[0]);
1830 break;
1831 }
1832
1833 case SpvOpSpecConstantComposite:
1834 case SpvOpConstantComposite: {
1835 unsigned elem_count = count - 3;
1836 vtn_fail_if(elem_count != val->type->length,
1837 "%s has %u constituents, expected %u",
1838 spirv_op_to_string(opcode), elem_count, val->type->length);
1839
1840 nir_constant **elems = ralloc_array(b, nir_constant *, elem_count);
1841 for (unsigned i = 0; i < elem_count; i++) {
1842 struct vtn_value *val = vtn_untyped_value(b, w[i + 3]);
1843
1844 if (val->value_type == vtn_value_type_constant) {
1845 elems[i] = val->constant;
1846 } else {
1847 vtn_fail_if(val->value_type != vtn_value_type_undef,
1848 "only constants or undefs allowed for "
1849 "SpvOpConstantComposite");
1850 /* to make it easier, just insert a NULL constant for now */
1851 elems[i] = vtn_null_constant(b, val->type);
1852 }
1853 }
1854
1855 switch (val->type->base_type) {
1856 case vtn_base_type_vector: {
1857 assert(glsl_type_is_vector(val->type->type));
1858 for (unsigned i = 0; i < elem_count; i++)
1859 val->constant->values[i] = elems[i]->values[0];
1860 break;
1861 }
1862
1863 case vtn_base_type_matrix:
1864 case vtn_base_type_struct:
1865 case vtn_base_type_array:
1866 ralloc_steal(val->constant, elems);
1867 val->constant->num_elements = elem_count;
1868 val->constant->elements = elems;
1869 break;
1870
1871 default:
1872 vtn_fail("Result type of %s must be a composite type",
1873 spirv_op_to_string(opcode));
1874 }
1875 break;
1876 }
1877
1878 case SpvOpSpecConstantOp: {
1879 nir_const_value u32op = nir_const_value_for_uint(w[3], 32);
1880 vtn_foreach_decoration(b, val, spec_constant_decoration_cb, &u32op);
1881 SpvOp opcode = u32op.u32;
1882 switch (opcode) {
1883 case SpvOpVectorShuffle: {
1884 struct vtn_value *v0 = &b->values[w[4]];
1885 struct vtn_value *v1 = &b->values[w[5]];
1886
1887 vtn_assert(v0->value_type == vtn_value_type_constant ||
1888 v0->value_type == vtn_value_type_undef);
1889 vtn_assert(v1->value_type == vtn_value_type_constant ||
1890 v1->value_type == vtn_value_type_undef);
1891
1892 unsigned len0 = glsl_get_vector_elements(v0->type->type);
1893 unsigned len1 = glsl_get_vector_elements(v1->type->type);
1894
1895 vtn_assert(len0 + len1 < 16);
1896
1897 unsigned bit_size = glsl_get_bit_size(val->type->type);
1898 unsigned bit_size0 = glsl_get_bit_size(v0->type->type);
1899 unsigned bit_size1 = glsl_get_bit_size(v1->type->type);
1900
1901 vtn_assert(bit_size == bit_size0 && bit_size == bit_size1);
1902 (void)bit_size0; (void)bit_size1;
1903
1904 nir_const_value undef = { .u64 = 0xdeadbeefdeadbeef };
1905 nir_const_value combined[NIR_MAX_VEC_COMPONENTS * 2];
1906
1907 if (v0->value_type == vtn_value_type_constant) {
1908 for (unsigned i = 0; i < len0; i++)
1909 combined[i] = v0->constant->values[i];
1910 }
1911 if (v1->value_type == vtn_value_type_constant) {
1912 for (unsigned i = 0; i < len1; i++)
1913 combined[len0 + i] = v1->constant->values[i];
1914 }
1915
1916 for (unsigned i = 0, j = 0; i < count - 6; i++, j++) {
1917 uint32_t comp = w[i + 6];
1918 if (comp == (uint32_t)-1) {
1919 /* If component is not used, set the value to a known constant
1920 * to detect if it is wrongly used.
1921 */
1922 val->constant->values[j] = undef;
1923 } else {
1924 vtn_fail_if(comp >= len0 + len1,
1925 "All Component literals must either be FFFFFFFF "
1926 "or in [0, N - 1] (inclusive).");
1927 val->constant->values[j] = combined[comp];
1928 }
1929 }
1930 break;
1931 }
1932
1933 case SpvOpCompositeExtract:
1934 case SpvOpCompositeInsert: {
1935 struct vtn_value *comp;
1936 unsigned deref_start;
1937 struct nir_constant **c;
1938 if (opcode == SpvOpCompositeExtract) {
1939 comp = vtn_value(b, w[4], vtn_value_type_constant);
1940 deref_start = 5;
1941 c = &comp->constant;
1942 } else {
1943 comp = vtn_value(b, w[5], vtn_value_type_constant);
1944 deref_start = 6;
1945 val->constant = nir_constant_clone(comp->constant,
1946 (nir_variable *)b);
1947 c = &val->constant;
1948 }
1949
1950 int elem = -1;
1951 const struct vtn_type *type = comp->type;
1952 for (unsigned i = deref_start; i < count; i++) {
1953 vtn_fail_if(w[i] > type->length,
1954 "%uth index of %s is %u but the type has only "
1955 "%u elements", i - deref_start,
1956 spirv_op_to_string(opcode), w[i], type->length);
1957
1958 switch (type->base_type) {
1959 case vtn_base_type_vector:
1960 elem = w[i];
1961 type = type->array_element;
1962 break;
1963
1964 case vtn_base_type_matrix:
1965 case vtn_base_type_array:
1966 c = &(*c)->elements[w[i]];
1967 type = type->array_element;
1968 break;
1969
1970 case vtn_base_type_struct:
1971 c = &(*c)->elements[w[i]];
1972 type = type->members[w[i]];
1973 break;
1974
1975 default:
1976 vtn_fail("%s must only index into composite types",
1977 spirv_op_to_string(opcode));
1978 }
1979 }
1980
1981 if (opcode == SpvOpCompositeExtract) {
1982 if (elem == -1) {
1983 val->constant = *c;
1984 } else {
1985 unsigned num_components = type->length;
1986 for (unsigned i = 0; i < num_components; i++)
1987 val->constant->values[i] = (*c)->values[elem + i];
1988 }
1989 } else {
1990 struct vtn_value *insert =
1991 vtn_value(b, w[4], vtn_value_type_constant);
1992 vtn_assert(insert->type == type);
1993 if (elem == -1) {
1994 *c = insert->constant;
1995 } else {
1996 unsigned num_components = type->length;
1997 for (unsigned i = 0; i < num_components; i++)
1998 (*c)->values[elem + i] = insert->constant->values[i];
1999 }
2000 }
2001 break;
2002 }
2003
2004 default: {
2005 bool swap;
2006 nir_alu_type dst_alu_type = nir_get_nir_type_for_glsl_type(val->type->type);
2007 nir_alu_type src_alu_type = dst_alu_type;
2008 unsigned num_components = glsl_get_vector_elements(val->type->type);
2009 unsigned bit_size;
2010
2011 vtn_assert(count <= 7);
2012
2013 switch (opcode) {
2014 case SpvOpSConvert:
2015 case SpvOpFConvert:
2016 case SpvOpUConvert:
2017 /* We have a source in a conversion */
2018 src_alu_type =
2019 nir_get_nir_type_for_glsl_type(vtn_get_value_type(b, w[4])->type);
2020 /* We use the bitsize of the conversion source to evaluate the opcode later */
2021 bit_size = glsl_get_bit_size(vtn_get_value_type(b, w[4])->type);
2022 break;
2023 default:
2024 bit_size = glsl_get_bit_size(val->type->type);
2025 };
2026
2027 nir_op op = vtn_nir_alu_op_for_spirv_opcode(b, opcode, &swap,
2028 nir_alu_type_get_type_size(src_alu_type),
2029 nir_alu_type_get_type_size(dst_alu_type));
2030 nir_const_value src[3][NIR_MAX_VEC_COMPONENTS];
2031
2032 for (unsigned i = 0; i < count - 4; i++) {
2033 struct vtn_value *src_val =
2034 vtn_value(b, w[4 + i], vtn_value_type_constant);
2035
2036 /* If this is an unsized source, pull the bit size from the
2037 * source; otherwise, we'll use the bit size from the destination.
2038 */
2039 if (!nir_alu_type_get_type_size(nir_op_infos[op].input_types[i]))
2040 bit_size = glsl_get_bit_size(src_val->type->type);
2041
2042 unsigned src_comps = nir_op_infos[op].input_sizes[i] ?
2043 nir_op_infos[op].input_sizes[i] :
2044 num_components;
2045
2046 unsigned j = swap ? 1 - i : i;
2047 for (unsigned c = 0; c < src_comps; c++)
2048 src[j][c] = src_val->constant->values[c];
2049 }
2050
2051 /* fix up fixed size sources */
2052 switch (op) {
2053 case nir_op_ishl:
2054 case nir_op_ishr:
2055 case nir_op_ushr: {
2056 if (bit_size == 32)
2057 break;
2058 for (unsigned i = 0; i < num_components; ++i) {
2059 switch (bit_size) {
2060 case 64: src[1][i].u32 = src[1][i].u64; break;
2061 case 16: src[1][i].u32 = src[1][i].u16; break;
2062 case 8: src[1][i].u32 = src[1][i].u8; break;
2063 }
2064 }
2065 break;
2066 }
2067 default:
2068 break;
2069 }
2070
2071 nir_const_value *srcs[3] = {
2072 src[0], src[1], src[2],
2073 };
2074 nir_eval_const_opcode(op, val->constant->values,
2075 num_components, bit_size, srcs,
2076 b->shader->info.float_controls_execution_mode);
2077 break;
2078 } /* default */
2079 }
2080 break;
2081 }
2082
2083 case SpvOpConstantNull:
2084 val->constant = vtn_null_constant(b, val->type);
2085 break;
2086
2087 case SpvOpConstantSampler:
2088 vtn_fail("OpConstantSampler requires Kernel Capability");
2089 break;
2090
2091 default:
2092 vtn_fail_with_opcode("Unhandled opcode", opcode);
2093 }
2094
2095 /* Now that we have the value, update the workgroup size if needed */
2096 vtn_foreach_decoration(b, val, handle_workgroup_size_decoration_cb, NULL);
2097 }
2098
2099 SpvMemorySemanticsMask
2100 vtn_storage_class_to_memory_semantics(SpvStorageClass sc)
2101 {
2102 switch (sc) {
2103 case SpvStorageClassStorageBuffer:
2104 case SpvStorageClassPhysicalStorageBuffer:
2105 return SpvMemorySemanticsUniformMemoryMask;
2106 case SpvStorageClassWorkgroup:
2107 return SpvMemorySemanticsWorkgroupMemoryMask;
2108 default:
2109 return SpvMemorySemanticsMaskNone;
2110 }
2111 }
2112
2113 static void
2114 vtn_split_barrier_semantics(struct vtn_builder *b,
2115 SpvMemorySemanticsMask semantics,
2116 SpvMemorySemanticsMask *before,
2117 SpvMemorySemanticsMask *after)
2118 {
2119 /* For memory semantics embedded in operations, we split them into up to
2120 * two barriers, to be added before and after the operation. This is less
2121 * strict than if we propagated until the final backend stage, but still
2122 * result in correct execution.
2123 *
2124 * A further improvement could be pipe this information (and use!) into the
2125 * next compiler layers, at the expense of making the handling of barriers
2126 * more complicated.
2127 */
2128
2129 *before = SpvMemorySemanticsMaskNone;
2130 *after = SpvMemorySemanticsMaskNone;
2131
2132 SpvMemorySemanticsMask order_semantics =
2133 semantics & (SpvMemorySemanticsAcquireMask |
2134 SpvMemorySemanticsReleaseMask |
2135 SpvMemorySemanticsAcquireReleaseMask |
2136 SpvMemorySemanticsSequentiallyConsistentMask);
2137
2138 if (util_bitcount(order_semantics) > 1) {
2139 /* Old GLSLang versions incorrectly set all the ordering bits. This was
2140 * fixed in c51287d744fb6e7e9ccc09f6f8451e6c64b1dad6 of glslang repo,
2141 * and it is in GLSLang since revision "SPIRV99.1321" (from Jul-2016).
2142 */
2143 vtn_warn("Multiple memory ordering semantics specified, "
2144 "assuming AcquireRelease.");
2145 order_semantics = SpvMemorySemanticsAcquireReleaseMask;
2146 }
2147
2148 const SpvMemorySemanticsMask av_vis_semantics =
2149 semantics & (SpvMemorySemanticsMakeAvailableMask |
2150 SpvMemorySemanticsMakeVisibleMask);
2151
2152 const SpvMemorySemanticsMask storage_semantics =
2153 semantics & (SpvMemorySemanticsUniformMemoryMask |
2154 SpvMemorySemanticsSubgroupMemoryMask |
2155 SpvMemorySemanticsWorkgroupMemoryMask |
2156 SpvMemorySemanticsCrossWorkgroupMemoryMask |
2157 SpvMemorySemanticsAtomicCounterMemoryMask |
2158 SpvMemorySemanticsImageMemoryMask |
2159 SpvMemorySemanticsOutputMemoryMask);
2160
2161 const SpvMemorySemanticsMask other_semantics =
2162 semantics & ~(order_semantics | av_vis_semantics | storage_semantics);
2163
2164 if (other_semantics)
2165 vtn_warn("Ignoring unhandled memory semantics: %u\n", other_semantics);
2166
2167 /* SequentiallyConsistent is treated as AcquireRelease. */
2168
2169 /* The RELEASE barrier happens BEFORE the operation, and it is usually
2170 * associated with a Store. All the write operations with a matching
2171 * semantics will not be reordered after the Store.
2172 */
2173 if (order_semantics & (SpvMemorySemanticsReleaseMask |
2174 SpvMemorySemanticsAcquireReleaseMask |
2175 SpvMemorySemanticsSequentiallyConsistentMask)) {
2176 *before |= SpvMemorySemanticsReleaseMask | storage_semantics;
2177 }
2178
2179 /* The ACQUIRE barrier happens AFTER the operation, and it is usually
2180 * associated with a Load. All the operations with a matching semantics
2181 * will not be reordered before the Load.
2182 */
2183 if (order_semantics & (SpvMemorySemanticsAcquireMask |
2184 SpvMemorySemanticsAcquireReleaseMask |
2185 SpvMemorySemanticsSequentiallyConsistentMask)) {
2186 *after |= SpvMemorySemanticsAcquireMask | storage_semantics;
2187 }
2188
2189 if (av_vis_semantics & SpvMemorySemanticsMakeVisibleMask)
2190 *before |= SpvMemorySemanticsMakeVisibleMask | storage_semantics;
2191
2192 if (av_vis_semantics & SpvMemorySemanticsMakeAvailableMask)
2193 *after |= SpvMemorySemanticsMakeAvailableMask | storage_semantics;
2194 }
2195
2196 static nir_memory_semantics
2197 vtn_mem_semantics_to_nir_mem_semantics(struct vtn_builder *b,
2198 SpvMemorySemanticsMask semantics)
2199 {
2200 nir_memory_semantics nir_semantics = 0;
2201
2202 SpvMemorySemanticsMask order_semantics =
2203 semantics & (SpvMemorySemanticsAcquireMask |
2204 SpvMemorySemanticsReleaseMask |
2205 SpvMemorySemanticsAcquireReleaseMask |
2206 SpvMemorySemanticsSequentiallyConsistentMask);
2207
2208 if (util_bitcount(order_semantics) > 1) {
2209 /* Old GLSLang versions incorrectly set all the ordering bits. This was
2210 * fixed in c51287d744fb6e7e9ccc09f6f8451e6c64b1dad6 of glslang repo,
2211 * and it is in GLSLang since revision "SPIRV99.1321" (from Jul-2016).
2212 */
2213 vtn_warn("Multiple memory ordering semantics bits specified, "
2214 "assuming AcquireRelease.");
2215 order_semantics = SpvMemorySemanticsAcquireReleaseMask;
2216 }
2217
2218 switch (order_semantics) {
2219 case 0:
2220 /* Not an ordering barrier. */
2221 break;
2222
2223 case SpvMemorySemanticsAcquireMask:
2224 nir_semantics = NIR_MEMORY_ACQUIRE;
2225 break;
2226
2227 case SpvMemorySemanticsReleaseMask:
2228 nir_semantics = NIR_MEMORY_RELEASE;
2229 break;
2230
2231 case SpvMemorySemanticsSequentiallyConsistentMask:
2232 /* Fall through. Treated as AcquireRelease in Vulkan. */
2233 case SpvMemorySemanticsAcquireReleaseMask:
2234 nir_semantics = NIR_MEMORY_ACQUIRE | NIR_MEMORY_RELEASE;
2235 break;
2236
2237 default:
2238 unreachable("Invalid memory order semantics");
2239 }
2240
2241 if (semantics & SpvMemorySemanticsMakeAvailableMask) {
2242 vtn_fail_if(!b->options->caps.vk_memory_model,
2243 "To use MakeAvailable memory semantics the VulkanMemoryModel "
2244 "capability must be declared.");
2245 nir_semantics |= NIR_MEMORY_MAKE_AVAILABLE;
2246 }
2247
2248 if (semantics & SpvMemorySemanticsMakeVisibleMask) {
2249 vtn_fail_if(!b->options->caps.vk_memory_model,
2250 "To use MakeVisible memory semantics the VulkanMemoryModel "
2251 "capability must be declared.");
2252 nir_semantics |= NIR_MEMORY_MAKE_VISIBLE;
2253 }
2254
2255 return nir_semantics;
2256 }
2257
2258 static nir_variable_mode
2259 vtn_mem_sematics_to_nir_var_modes(struct vtn_builder *b,
2260 SpvMemorySemanticsMask semantics)
2261 {
2262 /* Vulkan Environment for SPIR-V says "SubgroupMemory, CrossWorkgroupMemory,
2263 * and AtomicCounterMemory are ignored".
2264 */
2265 semantics &= ~(SpvMemorySemanticsSubgroupMemoryMask |
2266 SpvMemorySemanticsCrossWorkgroupMemoryMask |
2267 SpvMemorySemanticsAtomicCounterMemoryMask);
2268
2269 /* TODO: Consider adding nir_var_mem_image mode to NIR so it can be used
2270 * for SpvMemorySemanticsImageMemoryMask.
2271 */
2272
2273 nir_variable_mode modes = 0;
2274 if (semantics & (SpvMemorySemanticsUniformMemoryMask |
2275 SpvMemorySemanticsImageMemoryMask)) {
2276 modes |= nir_var_uniform |
2277 nir_var_mem_ubo |
2278 nir_var_mem_ssbo |
2279 nir_var_mem_global;
2280 }
2281 if (semantics & SpvMemorySemanticsWorkgroupMemoryMask)
2282 modes |= nir_var_mem_shared;
2283 if (semantics & SpvMemorySemanticsOutputMemoryMask) {
2284 modes |= nir_var_shader_out;
2285 }
2286
2287 return modes;
2288 }
2289
2290 static nir_scope
2291 vtn_scope_to_nir_scope(struct vtn_builder *b, SpvScope scope)
2292 {
2293 nir_scope nir_scope;
2294 switch (scope) {
2295 case SpvScopeDevice:
2296 vtn_fail_if(b->options->caps.vk_memory_model &&
2297 !b->options->caps.vk_memory_model_device_scope,
2298 "If the Vulkan memory model is declared and any instruction "
2299 "uses Device scope, the VulkanMemoryModelDeviceScope "
2300 "capability must be declared.");
2301 nir_scope = NIR_SCOPE_DEVICE;
2302 break;
2303
2304 case SpvScopeQueueFamily:
2305 vtn_fail_if(!b->options->caps.vk_memory_model,
2306 "To use Queue Family scope, the VulkanMemoryModel capability "
2307 "must be declared.");
2308 nir_scope = NIR_SCOPE_QUEUE_FAMILY;
2309 break;
2310
2311 case SpvScopeWorkgroup:
2312 nir_scope = NIR_SCOPE_WORKGROUP;
2313 break;
2314
2315 case SpvScopeSubgroup:
2316 nir_scope = NIR_SCOPE_SUBGROUP;
2317 break;
2318
2319 case SpvScopeInvocation:
2320 nir_scope = NIR_SCOPE_INVOCATION;
2321 break;
2322
2323 default:
2324 vtn_fail("Invalid memory scope");
2325 }
2326
2327 return nir_scope;
2328 }
2329
2330 static void
2331 vtn_emit_scoped_control_barrier(struct vtn_builder *b, SpvScope exec_scope,
2332 SpvScope mem_scope,
2333 SpvMemorySemanticsMask semantics)
2334 {
2335 nir_memory_semantics nir_semantics =
2336 vtn_mem_semantics_to_nir_mem_semantics(b, semantics);
2337 nir_variable_mode modes = vtn_mem_sematics_to_nir_var_modes(b, semantics);
2338 nir_scope nir_exec_scope = vtn_scope_to_nir_scope(b, exec_scope);
2339
2340 /* Memory semantics is optional for OpControlBarrier. */
2341 nir_scope nir_mem_scope;
2342 if (nir_semantics == 0 || modes == 0)
2343 nir_mem_scope = NIR_SCOPE_NONE;
2344 else
2345 nir_mem_scope = vtn_scope_to_nir_scope(b, mem_scope);
2346
2347 nir_scoped_barrier(&b->nb, nir_exec_scope, nir_mem_scope, nir_semantics, modes);
2348 }
2349
2350 static void
2351 vtn_emit_scoped_memory_barrier(struct vtn_builder *b, SpvScope scope,
2352 SpvMemorySemanticsMask semantics)
2353 {
2354 nir_variable_mode modes = vtn_mem_sematics_to_nir_var_modes(b, semantics);
2355 nir_memory_semantics nir_semantics =
2356 vtn_mem_semantics_to_nir_mem_semantics(b, semantics);
2357
2358 /* No barrier to add. */
2359 if (nir_semantics == 0 || modes == 0)
2360 return;
2361
2362 nir_scope nir_mem_scope = vtn_scope_to_nir_scope(b, scope);
2363 nir_scoped_barrier(&b->nb, NIR_SCOPE_NONE, nir_mem_scope, nir_semantics, modes);
2364 }
2365
2366 struct vtn_ssa_value *
2367 vtn_create_ssa_value(struct vtn_builder *b, const struct glsl_type *type)
2368 {
2369 /* Always use bare types for SSA values for a couple of reasons:
2370 *
2371 * 1. Code which emits deref chains should never listen to the explicit
2372 * layout information on the SSA value if any exists. If we've
2373 * accidentally been relying on this, we want to find those bugs.
2374 *
2375 * 2. We want to be able to quickly check that an SSA value being assigned
2376 * to a SPIR-V value has the right type. Using bare types everywhere
2377 * ensures that we can pointer-compare.
2378 */
2379 struct vtn_ssa_value *val = rzalloc(b, struct vtn_ssa_value);
2380 val->type = glsl_get_bare_type(type);
2381
2382
2383 if (!glsl_type_is_vector_or_scalar(type)) {
2384 unsigned elems = glsl_get_length(val->type);
2385 val->elems = ralloc_array(b, struct vtn_ssa_value *, elems);
2386 if (glsl_type_is_array_or_matrix(type)) {
2387 const struct glsl_type *elem_type = glsl_get_array_element(type);
2388 for (unsigned i = 0; i < elems; i++)
2389 val->elems[i] = vtn_create_ssa_value(b, elem_type);
2390 } else {
2391 vtn_assert(glsl_type_is_struct_or_ifc(type));
2392 for (unsigned i = 0; i < elems; i++) {
2393 const struct glsl_type *elem_type = glsl_get_struct_field(type, i);
2394 val->elems[i] = vtn_create_ssa_value(b, elem_type);
2395 }
2396 }
2397 }
2398
2399 return val;
2400 }
2401
2402 static nir_tex_src
2403 vtn_tex_src(struct vtn_builder *b, unsigned index, nir_tex_src_type type)
2404 {
2405 nir_tex_src src;
2406 src.src = nir_src_for_ssa(vtn_get_nir_ssa(b, index));
2407 src.src_type = type;
2408 return src;
2409 }
2410
2411 static uint32_t
2412 image_operand_arg(struct vtn_builder *b, const uint32_t *w, uint32_t count,
2413 uint32_t mask_idx, SpvImageOperandsMask op)
2414 {
2415 static const SpvImageOperandsMask ops_with_arg =
2416 SpvImageOperandsBiasMask |
2417 SpvImageOperandsLodMask |
2418 SpvImageOperandsGradMask |
2419 SpvImageOperandsConstOffsetMask |
2420 SpvImageOperandsOffsetMask |
2421 SpvImageOperandsConstOffsetsMask |
2422 SpvImageOperandsSampleMask |
2423 SpvImageOperandsMinLodMask |
2424 SpvImageOperandsMakeTexelAvailableMask |
2425 SpvImageOperandsMakeTexelVisibleMask;
2426
2427 assert(util_bitcount(op) == 1);
2428 assert(w[mask_idx] & op);
2429 assert(op & ops_with_arg);
2430
2431 uint32_t idx = util_bitcount(w[mask_idx] & (op - 1) & ops_with_arg) + 1;
2432
2433 /* Adjust indices for operands with two arguments. */
2434 static const SpvImageOperandsMask ops_with_two_args =
2435 SpvImageOperandsGradMask;
2436 idx += util_bitcount(w[mask_idx] & (op - 1) & ops_with_two_args);
2437
2438 idx += mask_idx;
2439
2440 vtn_fail_if(idx + (op & ops_with_two_args ? 1 : 0) >= count,
2441 "Image op claims to have %s but does not enough "
2442 "following operands", spirv_imageoperands_to_string(op));
2443
2444 return idx;
2445 }
2446
2447 static void
2448 non_uniform_decoration_cb(struct vtn_builder *b,
2449 struct vtn_value *val, int member,
2450 const struct vtn_decoration *dec, void *void_ctx)
2451 {
2452 enum gl_access_qualifier *access = void_ctx;
2453 switch (dec->decoration) {
2454 case SpvDecorationNonUniformEXT:
2455 *access |= ACCESS_NON_UNIFORM;
2456 break;
2457
2458 default:
2459 break;
2460 }
2461 }
2462
2463 static void
2464 vtn_handle_texture(struct vtn_builder *b, SpvOp opcode,
2465 const uint32_t *w, unsigned count)
2466 {
2467 struct vtn_type *ret_type = vtn_get_type(b, w[1]);
2468
2469 if (opcode == SpvOpSampledImage) {
2470 struct vtn_sampled_image si = {
2471 .image = vtn_get_image(b, w[3]),
2472 .sampler = vtn_get_sampler(b, w[4]),
2473 };
2474 vtn_push_sampled_image(b, w[2], si);
2475 return;
2476 } else if (opcode == SpvOpImage) {
2477 struct vtn_sampled_image si = vtn_get_sampled_image(b, w[3]);
2478 vtn_push_image(b, w[2], si.image);
2479 return;
2480 }
2481
2482 nir_deref_instr *image = NULL, *sampler = NULL;
2483 struct vtn_value *sampled_val = vtn_untyped_value(b, w[3]);
2484 if (sampled_val->type->base_type == vtn_base_type_sampled_image) {
2485 struct vtn_sampled_image si = vtn_get_sampled_image(b, w[3]);
2486 image = si.image;
2487 sampler = si.sampler;
2488 } else {
2489 image = vtn_get_image(b, w[3]);
2490 }
2491
2492 const enum glsl_sampler_dim sampler_dim = glsl_get_sampler_dim(image->type);
2493 const bool is_array = glsl_sampler_type_is_array(image->type);
2494 nir_alu_type dest_type = nir_type_invalid;
2495
2496 /* Figure out the base texture operation */
2497 nir_texop texop;
2498 switch (opcode) {
2499 case SpvOpImageSampleImplicitLod:
2500 case SpvOpImageSampleDrefImplicitLod:
2501 case SpvOpImageSampleProjImplicitLod:
2502 case SpvOpImageSampleProjDrefImplicitLod:
2503 texop = nir_texop_tex;
2504 break;
2505
2506 case SpvOpImageSampleExplicitLod:
2507 case SpvOpImageSampleDrefExplicitLod:
2508 case SpvOpImageSampleProjExplicitLod:
2509 case SpvOpImageSampleProjDrefExplicitLod:
2510 texop = nir_texop_txl;
2511 break;
2512
2513 case SpvOpImageFetch:
2514 if (sampler_dim == GLSL_SAMPLER_DIM_MS) {
2515 texop = nir_texop_txf_ms;
2516 } else {
2517 texop = nir_texop_txf;
2518 }
2519 break;
2520
2521 case SpvOpImageGather:
2522 case SpvOpImageDrefGather:
2523 texop = nir_texop_tg4;
2524 break;
2525
2526 case SpvOpImageQuerySizeLod:
2527 case SpvOpImageQuerySize:
2528 texop = nir_texop_txs;
2529 dest_type = nir_type_int;
2530 break;
2531
2532 case SpvOpImageQueryLod:
2533 texop = nir_texop_lod;
2534 dest_type = nir_type_float;
2535 break;
2536
2537 case SpvOpImageQueryLevels:
2538 texop = nir_texop_query_levels;
2539 dest_type = nir_type_int;
2540 break;
2541
2542 case SpvOpImageQuerySamples:
2543 texop = nir_texop_texture_samples;
2544 dest_type = nir_type_int;
2545 break;
2546
2547 case SpvOpFragmentFetchAMD:
2548 texop = nir_texop_fragment_fetch;
2549 break;
2550
2551 case SpvOpFragmentMaskFetchAMD:
2552 texop = nir_texop_fragment_mask_fetch;
2553 break;
2554
2555 default:
2556 vtn_fail_with_opcode("Unhandled opcode", opcode);
2557 }
2558
2559 nir_tex_src srcs[10]; /* 10 should be enough */
2560 nir_tex_src *p = srcs;
2561
2562 p->src = nir_src_for_ssa(&image->dest.ssa);
2563 p->src_type = nir_tex_src_texture_deref;
2564 p++;
2565
2566 switch (texop) {
2567 case nir_texop_tex:
2568 case nir_texop_txb:
2569 case nir_texop_txl:
2570 case nir_texop_txd:
2571 case nir_texop_tg4:
2572 case nir_texop_lod:
2573 vtn_fail_if(sampler == NULL,
2574 "%s requires an image of type OpTypeSampledImage",
2575 spirv_op_to_string(opcode));
2576 p->src = nir_src_for_ssa(&sampler->dest.ssa);
2577 p->src_type = nir_tex_src_sampler_deref;
2578 p++;
2579 break;
2580 case nir_texop_txf:
2581 case nir_texop_txf_ms:
2582 case nir_texop_txs:
2583 case nir_texop_query_levels:
2584 case nir_texop_texture_samples:
2585 case nir_texop_samples_identical:
2586 case nir_texop_fragment_fetch:
2587 case nir_texop_fragment_mask_fetch:
2588 /* These don't */
2589 break;
2590 case nir_texop_txf_ms_fb:
2591 vtn_fail("unexpected nir_texop_txf_ms_fb");
2592 break;
2593 case nir_texop_txf_ms_mcs:
2594 vtn_fail("unexpected nir_texop_txf_ms_mcs");
2595 case nir_texop_tex_prefetch:
2596 vtn_fail("unexpected nir_texop_tex_prefetch");
2597 }
2598
2599 unsigned idx = 4;
2600
2601 struct nir_ssa_def *coord;
2602 unsigned coord_components;
2603 switch (opcode) {
2604 case SpvOpImageSampleImplicitLod:
2605 case SpvOpImageSampleExplicitLod:
2606 case SpvOpImageSampleDrefImplicitLod:
2607 case SpvOpImageSampleDrefExplicitLod:
2608 case SpvOpImageSampleProjImplicitLod:
2609 case SpvOpImageSampleProjExplicitLod:
2610 case SpvOpImageSampleProjDrefImplicitLod:
2611 case SpvOpImageSampleProjDrefExplicitLod:
2612 case SpvOpImageFetch:
2613 case SpvOpImageGather:
2614 case SpvOpImageDrefGather:
2615 case SpvOpImageQueryLod:
2616 case SpvOpFragmentFetchAMD:
2617 case SpvOpFragmentMaskFetchAMD: {
2618 /* All these types have the coordinate as their first real argument */
2619 coord_components = glsl_get_sampler_dim_coordinate_components(sampler_dim);
2620
2621 if (is_array && texop != nir_texop_lod)
2622 coord_components++;
2623
2624 coord = vtn_get_nir_ssa(b, w[idx++]);
2625 p->src = nir_src_for_ssa(nir_channels(&b->nb, coord,
2626 (1 << coord_components) - 1));
2627 p->src_type = nir_tex_src_coord;
2628 p++;
2629 break;
2630 }
2631
2632 default:
2633 coord = NULL;
2634 coord_components = 0;
2635 break;
2636 }
2637
2638 switch (opcode) {
2639 case SpvOpImageSampleProjImplicitLod:
2640 case SpvOpImageSampleProjExplicitLod:
2641 case SpvOpImageSampleProjDrefImplicitLod:
2642 case SpvOpImageSampleProjDrefExplicitLod:
2643 /* These have the projector as the last coordinate component */
2644 p->src = nir_src_for_ssa(nir_channel(&b->nb, coord, coord_components));
2645 p->src_type = nir_tex_src_projector;
2646 p++;
2647 break;
2648
2649 default:
2650 break;
2651 }
2652
2653 bool is_shadow = false;
2654 unsigned gather_component = 0;
2655 switch (opcode) {
2656 case SpvOpImageSampleDrefImplicitLod:
2657 case SpvOpImageSampleDrefExplicitLod:
2658 case SpvOpImageSampleProjDrefImplicitLod:
2659 case SpvOpImageSampleProjDrefExplicitLod:
2660 case SpvOpImageDrefGather:
2661 /* These all have an explicit depth value as their next source */
2662 is_shadow = true;
2663 (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_comparator);
2664 break;
2665
2666 case SpvOpImageGather:
2667 /* This has a component as its next source */
2668 gather_component = vtn_constant_uint(b, w[idx++]);
2669 break;
2670
2671 default:
2672 break;
2673 }
2674
2675 /* For OpImageQuerySizeLod, we always have an LOD */
2676 if (opcode == SpvOpImageQuerySizeLod)
2677 (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_lod);
2678
2679 /* For OpFragmentFetchAMD, we always have a multisample index */
2680 if (opcode == SpvOpFragmentFetchAMD)
2681 (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_ms_index);
2682
2683 /* Now we need to handle some number of optional arguments */
2684 struct vtn_value *gather_offsets = NULL;
2685 if (idx < count) {
2686 uint32_t operands = w[idx];
2687
2688 if (operands & SpvImageOperandsBiasMask) {
2689 vtn_assert(texop == nir_texop_tex ||
2690 texop == nir_texop_tg4);
2691 if (texop == nir_texop_tex)
2692 texop = nir_texop_txb;
2693 uint32_t arg = image_operand_arg(b, w, count, idx,
2694 SpvImageOperandsBiasMask);
2695 (*p++) = vtn_tex_src(b, w[arg], nir_tex_src_bias);
2696 }
2697
2698 if (operands & SpvImageOperandsLodMask) {
2699 vtn_assert(texop == nir_texop_txl || texop == nir_texop_txf ||
2700 texop == nir_texop_txs || texop == nir_texop_tg4);
2701 uint32_t arg = image_operand_arg(b, w, count, idx,
2702 SpvImageOperandsLodMask);
2703 (*p++) = vtn_tex_src(b, w[arg], nir_tex_src_lod);
2704 }
2705
2706 if (operands & SpvImageOperandsGradMask) {
2707 vtn_assert(texop == nir_texop_txl);
2708 texop = nir_texop_txd;
2709 uint32_t arg = image_operand_arg(b, w, count, idx,
2710 SpvImageOperandsGradMask);
2711 (*p++) = vtn_tex_src(b, w[arg], nir_tex_src_ddx);
2712 (*p++) = vtn_tex_src(b, w[arg + 1], nir_tex_src_ddy);
2713 }
2714
2715 vtn_fail_if(util_bitcount(operands & (SpvImageOperandsConstOffsetsMask |
2716 SpvImageOperandsOffsetMask |
2717 SpvImageOperandsConstOffsetMask)) > 1,
2718 "At most one of the ConstOffset, Offset, and ConstOffsets "
2719 "image operands can be used on a given instruction.");
2720
2721 if (operands & SpvImageOperandsOffsetMask) {
2722 uint32_t arg = image_operand_arg(b, w, count, idx,
2723 SpvImageOperandsOffsetMask);
2724 (*p++) = vtn_tex_src(b, w[arg], nir_tex_src_offset);
2725 }
2726
2727 if (operands & SpvImageOperandsConstOffsetMask) {
2728 uint32_t arg = image_operand_arg(b, w, count, idx,
2729 SpvImageOperandsConstOffsetMask);
2730 (*p++) = vtn_tex_src(b, w[arg], nir_tex_src_offset);
2731 }
2732
2733 if (operands & SpvImageOperandsConstOffsetsMask) {
2734 vtn_assert(texop == nir_texop_tg4);
2735 uint32_t arg = image_operand_arg(b, w, count, idx,
2736 SpvImageOperandsConstOffsetsMask);
2737 gather_offsets = vtn_value(b, w[arg], vtn_value_type_constant);
2738 }
2739
2740 if (operands & SpvImageOperandsSampleMask) {
2741 vtn_assert(texop == nir_texop_txf_ms);
2742 uint32_t arg = image_operand_arg(b, w, count, idx,
2743 SpvImageOperandsSampleMask);
2744 texop = nir_texop_txf_ms;
2745 (*p++) = vtn_tex_src(b, w[arg], nir_tex_src_ms_index);
2746 }
2747
2748 if (operands & SpvImageOperandsMinLodMask) {
2749 vtn_assert(texop == nir_texop_tex ||
2750 texop == nir_texop_txb ||
2751 texop == nir_texop_txd);
2752 uint32_t arg = image_operand_arg(b, w, count, idx,
2753 SpvImageOperandsMinLodMask);
2754 (*p++) = vtn_tex_src(b, w[arg], nir_tex_src_min_lod);
2755 }
2756 }
2757
2758 nir_tex_instr *instr = nir_tex_instr_create(b->shader, p - srcs);
2759 instr->op = texop;
2760
2761 memcpy(instr->src, srcs, instr->num_srcs * sizeof(*instr->src));
2762
2763 instr->coord_components = coord_components;
2764 instr->sampler_dim = sampler_dim;
2765 instr->is_array = is_array;
2766 instr->is_shadow = is_shadow;
2767 instr->is_new_style_shadow =
2768 is_shadow && glsl_get_components(ret_type->type) == 1;
2769 instr->component = gather_component;
2770
2771 /* The Vulkan spec says:
2772 *
2773 * "If an instruction loads from or stores to a resource (including
2774 * atomics and image instructions) and the resource descriptor being
2775 * accessed is not dynamically uniform, then the operand corresponding
2776 * to that resource (e.g. the pointer or sampled image operand) must be
2777 * decorated with NonUniform."
2778 *
2779 * It's very careful to specify that the exact operand must be decorated
2780 * NonUniform. The SPIR-V parser is not expected to chase through long
2781 * chains to find the NonUniform decoration. It's either right there or we
2782 * can assume it doesn't exist.
2783 */
2784 enum gl_access_qualifier access = 0;
2785 vtn_foreach_decoration(b, sampled_val, non_uniform_decoration_cb, &access);
2786
2787 if (image && (access & ACCESS_NON_UNIFORM))
2788 instr->texture_non_uniform = true;
2789
2790 if (sampler && (access & ACCESS_NON_UNIFORM))
2791 instr->sampler_non_uniform = true;
2792
2793 /* for non-query ops, get dest_type from sampler type */
2794 if (dest_type == nir_type_invalid) {
2795 switch (glsl_get_sampler_result_type(image->type)) {
2796 case GLSL_TYPE_FLOAT: dest_type = nir_type_float; break;
2797 case GLSL_TYPE_INT: dest_type = nir_type_int; break;
2798 case GLSL_TYPE_UINT: dest_type = nir_type_uint; break;
2799 case GLSL_TYPE_BOOL: dest_type = nir_type_bool; break;
2800 default:
2801 vtn_fail("Invalid base type for sampler result");
2802 }
2803 }
2804
2805 instr->dest_type = dest_type;
2806
2807 nir_ssa_dest_init(&instr->instr, &instr->dest,
2808 nir_tex_instr_dest_size(instr), 32, NULL);
2809
2810 vtn_assert(glsl_get_vector_elements(ret_type->type) ==
2811 nir_tex_instr_dest_size(instr));
2812
2813 if (gather_offsets) {
2814 vtn_fail_if(gather_offsets->type->base_type != vtn_base_type_array ||
2815 gather_offsets->type->length != 4,
2816 "ConstOffsets must be an array of size four of vectors "
2817 "of two integer components");
2818
2819 struct vtn_type *vec_type = gather_offsets->type->array_element;
2820 vtn_fail_if(vec_type->base_type != vtn_base_type_vector ||
2821 vec_type->length != 2 ||
2822 !glsl_type_is_integer(vec_type->type),
2823 "ConstOffsets must be an array of size four of vectors "
2824 "of two integer components");
2825
2826 unsigned bit_size = glsl_get_bit_size(vec_type->type);
2827 for (uint32_t i = 0; i < 4; i++) {
2828 const nir_const_value *cvec =
2829 gather_offsets->constant->elements[i]->values;
2830 for (uint32_t j = 0; j < 2; j++) {
2831 switch (bit_size) {
2832 case 8: instr->tg4_offsets[i][j] = cvec[j].i8; break;
2833 case 16: instr->tg4_offsets[i][j] = cvec[j].i16; break;
2834 case 32: instr->tg4_offsets[i][j] = cvec[j].i32; break;
2835 case 64: instr->tg4_offsets[i][j] = cvec[j].i64; break;
2836 default:
2837 vtn_fail("Unsupported bit size: %u", bit_size);
2838 }
2839 }
2840 }
2841 }
2842
2843 nir_builder_instr_insert(&b->nb, &instr->instr);
2844
2845 vtn_push_nir_ssa(b, w[2], &instr->dest.ssa);
2846 }
2847
2848 static void
2849 fill_common_atomic_sources(struct vtn_builder *b, SpvOp opcode,
2850 const uint32_t *w, nir_src *src)
2851 {
2852 switch (opcode) {
2853 case SpvOpAtomicIIncrement:
2854 src[0] = nir_src_for_ssa(nir_imm_int(&b->nb, 1));
2855 break;
2856
2857 case SpvOpAtomicIDecrement:
2858 src[0] = nir_src_for_ssa(nir_imm_int(&b->nb, -1));
2859 break;
2860
2861 case SpvOpAtomicISub:
2862 src[0] =
2863 nir_src_for_ssa(nir_ineg(&b->nb, vtn_get_nir_ssa(b, w[6])));
2864 break;
2865
2866 case SpvOpAtomicCompareExchange:
2867 case SpvOpAtomicCompareExchangeWeak:
2868 src[0] = nir_src_for_ssa(vtn_get_nir_ssa(b, w[8]));
2869 src[1] = nir_src_for_ssa(vtn_get_nir_ssa(b, w[7]));
2870 break;
2871
2872 case SpvOpAtomicExchange:
2873 case SpvOpAtomicIAdd:
2874 case SpvOpAtomicSMin:
2875 case SpvOpAtomicUMin:
2876 case SpvOpAtomicSMax:
2877 case SpvOpAtomicUMax:
2878 case SpvOpAtomicAnd:
2879 case SpvOpAtomicOr:
2880 case SpvOpAtomicXor:
2881 case SpvOpAtomicFAddEXT:
2882 src[0] = nir_src_for_ssa(vtn_get_nir_ssa(b, w[6]));
2883 break;
2884
2885 default:
2886 vtn_fail_with_opcode("Invalid SPIR-V atomic", opcode);
2887 }
2888 }
2889
2890 static nir_ssa_def *
2891 get_image_coord(struct vtn_builder *b, uint32_t value)
2892 {
2893 nir_ssa_def *coord = vtn_get_nir_ssa(b, value);
2894
2895 /* The image_load_store intrinsics assume a 4-dim coordinate */
2896 unsigned swizzle[4];
2897 for (unsigned i = 0; i < 4; i++)
2898 swizzle[i] = MIN2(i, coord->num_components - 1);
2899
2900 return nir_swizzle(&b->nb, coord, swizzle, 4);
2901 }
2902
2903 static nir_ssa_def *
2904 expand_to_vec4(nir_builder *b, nir_ssa_def *value)
2905 {
2906 if (value->num_components == 4)
2907 return value;
2908
2909 unsigned swiz[4];
2910 for (unsigned i = 0; i < 4; i++)
2911 swiz[i] = i < value->num_components ? i : 0;
2912 return nir_swizzle(b, value, swiz, 4);
2913 }
2914
2915 static void
2916 vtn_handle_image(struct vtn_builder *b, SpvOp opcode,
2917 const uint32_t *w, unsigned count)
2918 {
2919 /* Just get this one out of the way */
2920 if (opcode == SpvOpImageTexelPointer) {
2921 struct vtn_value *val =
2922 vtn_push_value(b, w[2], vtn_value_type_image_pointer);
2923 val->image = ralloc(b, struct vtn_image_pointer);
2924
2925 val->image->image = vtn_nir_deref(b, w[3]);
2926 val->image->coord = get_image_coord(b, w[4]);
2927 val->image->sample = vtn_get_nir_ssa(b, w[5]);
2928 val->image->lod = nir_imm_int(&b->nb, 0);
2929 return;
2930 }
2931
2932 struct vtn_image_pointer image;
2933 SpvScope scope = SpvScopeInvocation;
2934 SpvMemorySemanticsMask semantics = 0;
2935
2936 enum gl_access_qualifier access = 0;
2937
2938 struct vtn_value *res_val;
2939 switch (opcode) {
2940 case SpvOpAtomicExchange:
2941 case SpvOpAtomicCompareExchange:
2942 case SpvOpAtomicCompareExchangeWeak:
2943 case SpvOpAtomicIIncrement:
2944 case SpvOpAtomicIDecrement:
2945 case SpvOpAtomicIAdd:
2946 case SpvOpAtomicISub:
2947 case SpvOpAtomicLoad:
2948 case SpvOpAtomicSMin:
2949 case SpvOpAtomicUMin:
2950 case SpvOpAtomicSMax:
2951 case SpvOpAtomicUMax:
2952 case SpvOpAtomicAnd:
2953 case SpvOpAtomicOr:
2954 case SpvOpAtomicXor:
2955 case SpvOpAtomicFAddEXT:
2956 res_val = vtn_value(b, w[3], vtn_value_type_image_pointer);
2957 image = *res_val->image;
2958 scope = vtn_constant_uint(b, w[4]);
2959 semantics = vtn_constant_uint(b, w[5]);
2960 access |= ACCESS_COHERENT;
2961 break;
2962
2963 case SpvOpAtomicStore:
2964 res_val = vtn_value(b, w[1], vtn_value_type_image_pointer);
2965 image = *res_val->image;
2966 scope = vtn_constant_uint(b, w[2]);
2967 semantics = vtn_constant_uint(b, w[3]);
2968 access |= ACCESS_COHERENT;
2969 break;
2970
2971 case SpvOpImageQuerySize:
2972 res_val = vtn_untyped_value(b, w[3]);
2973 image.image = vtn_get_image(b, w[3]);
2974 image.coord = NULL;
2975 image.sample = NULL;
2976 image.lod = NULL;
2977 break;
2978
2979 case SpvOpImageRead: {
2980 res_val = vtn_untyped_value(b, w[3]);
2981 image.image = vtn_get_image(b, w[3]);
2982 image.coord = get_image_coord(b, w[4]);
2983
2984 const SpvImageOperandsMask operands =
2985 count > 5 ? w[5] : SpvImageOperandsMaskNone;
2986
2987 if (operands & SpvImageOperandsSampleMask) {
2988 uint32_t arg = image_operand_arg(b, w, count, 5,
2989 SpvImageOperandsSampleMask);
2990 image.sample = vtn_get_nir_ssa(b, w[arg]);
2991 } else {
2992 image.sample = nir_ssa_undef(&b->nb, 1, 32);
2993 }
2994
2995 if (operands & SpvImageOperandsMakeTexelVisibleMask) {
2996 vtn_fail_if((operands & SpvImageOperandsNonPrivateTexelMask) == 0,
2997 "MakeTexelVisible requires NonPrivateTexel to also be set.");
2998 uint32_t arg = image_operand_arg(b, w, count, 5,
2999 SpvImageOperandsMakeTexelVisibleMask);
3000 semantics = SpvMemorySemanticsMakeVisibleMask;
3001 scope = vtn_constant_uint(b, w[arg]);
3002 }
3003
3004 if (operands & SpvImageOperandsLodMask) {
3005 uint32_t arg = image_operand_arg(b, w, count, 5,
3006 SpvImageOperandsLodMask);
3007 image.lod = vtn_get_nir_ssa(b, w[arg]);
3008 } else {
3009 image.lod = nir_imm_int(&b->nb, 0);
3010 }
3011
3012 /* TODO: Volatile. */
3013
3014 break;
3015 }
3016
3017 case SpvOpImageWrite: {
3018 res_val = vtn_untyped_value(b, w[1]);
3019 image.image = vtn_get_image(b, w[1]);
3020 image.coord = get_image_coord(b, w[2]);
3021
3022 /* texel = w[3] */
3023
3024 const SpvImageOperandsMask operands =
3025 count > 4 ? w[4] : SpvImageOperandsMaskNone;
3026
3027 if (operands & SpvImageOperandsSampleMask) {
3028 uint32_t arg = image_operand_arg(b, w, count, 4,
3029 SpvImageOperandsSampleMask);
3030 image.sample = vtn_get_nir_ssa(b, w[arg]);
3031 } else {
3032 image.sample = nir_ssa_undef(&b->nb, 1, 32);
3033 }
3034
3035 if (operands & SpvImageOperandsMakeTexelAvailableMask) {
3036 vtn_fail_if((operands & SpvImageOperandsNonPrivateTexelMask) == 0,
3037 "MakeTexelAvailable requires NonPrivateTexel to also be set.");
3038 uint32_t arg = image_operand_arg(b, w, count, 4,
3039 SpvImageOperandsMakeTexelAvailableMask);
3040 semantics = SpvMemorySemanticsMakeAvailableMask;
3041 scope = vtn_constant_uint(b, w[arg]);
3042 }
3043
3044 if (operands & SpvImageOperandsLodMask) {
3045 uint32_t arg = image_operand_arg(b, w, count, 4,
3046 SpvImageOperandsLodMask);
3047 image.lod = vtn_get_nir_ssa(b, w[arg]);
3048 } else {
3049 image.lod = nir_imm_int(&b->nb, 0);
3050 }
3051
3052 /* TODO: Volatile. */
3053
3054 break;
3055 }
3056
3057 default:
3058 vtn_fail_with_opcode("Invalid image opcode", opcode);
3059 }
3060
3061 nir_intrinsic_op op;
3062 switch (opcode) {
3063 #define OP(S, N) case SpvOp##S: op = nir_intrinsic_image_deref_##N; break;
3064 OP(ImageQuerySize, size)
3065 OP(ImageRead, load)
3066 OP(ImageWrite, store)
3067 OP(AtomicLoad, load)
3068 OP(AtomicStore, store)
3069 OP(AtomicExchange, atomic_exchange)
3070 OP(AtomicCompareExchange, atomic_comp_swap)
3071 OP(AtomicCompareExchangeWeak, atomic_comp_swap)
3072 OP(AtomicIIncrement, atomic_add)
3073 OP(AtomicIDecrement, atomic_add)
3074 OP(AtomicIAdd, atomic_add)
3075 OP(AtomicISub, atomic_add)
3076 OP(AtomicSMin, atomic_imin)
3077 OP(AtomicUMin, atomic_umin)
3078 OP(AtomicSMax, atomic_imax)
3079 OP(AtomicUMax, atomic_umax)
3080 OP(AtomicAnd, atomic_and)
3081 OP(AtomicOr, atomic_or)
3082 OP(AtomicXor, atomic_xor)
3083 OP(AtomicFAddEXT, atomic_fadd)
3084 #undef OP
3085 default:
3086 vtn_fail_with_opcode("Invalid image opcode", opcode);
3087 }
3088
3089 nir_intrinsic_instr *intrin = nir_intrinsic_instr_create(b->shader, op);
3090
3091 intrin->src[0] = nir_src_for_ssa(&image.image->dest.ssa);
3092
3093 if (opcode == SpvOpImageQuerySize) {
3094 /* ImageQuerySize only has an LOD which is currently always 0 */
3095 intrin->src[1] = nir_src_for_ssa(nir_imm_int(&b->nb, 0));
3096 } else {
3097 /* The image coordinate is always 4 components but we may not have that
3098 * many. Swizzle to compensate.
3099 */
3100 intrin->src[1] = nir_src_for_ssa(expand_to_vec4(&b->nb, image.coord));
3101 intrin->src[2] = nir_src_for_ssa(image.sample);
3102 }
3103
3104 /* The Vulkan spec says:
3105 *
3106 * "If an instruction loads from or stores to a resource (including
3107 * atomics and image instructions) and the resource descriptor being
3108 * accessed is not dynamically uniform, then the operand corresponding
3109 * to that resource (e.g. the pointer or sampled image operand) must be
3110 * decorated with NonUniform."
3111 *
3112 * It's very careful to specify that the exact operand must be decorated
3113 * NonUniform. The SPIR-V parser is not expected to chase through long
3114 * chains to find the NonUniform decoration. It's either right there or we
3115 * can assume it doesn't exist.
3116 */
3117 vtn_foreach_decoration(b, res_val, non_uniform_decoration_cb, &access);
3118 nir_intrinsic_set_access(intrin, access);
3119
3120 switch (opcode) {
3121 case SpvOpAtomicLoad:
3122 case SpvOpImageQuerySize:
3123 case SpvOpImageRead:
3124 if (opcode == SpvOpImageRead || opcode == SpvOpAtomicLoad) {
3125 /* Only OpImageRead can support a lod parameter if
3126 * SPV_AMD_shader_image_load_store_lod is used but the current NIR
3127 * intrinsics definition for atomics requires us to set it for
3128 * OpAtomicLoad.
3129 */
3130 intrin->src[3] = nir_src_for_ssa(image.lod);
3131 }
3132 break;
3133 case SpvOpAtomicStore:
3134 case SpvOpImageWrite: {
3135 const uint32_t value_id = opcode == SpvOpAtomicStore ? w[4] : w[3];
3136 nir_ssa_def *value = vtn_get_nir_ssa(b, value_id);
3137 /* nir_intrinsic_image_deref_store always takes a vec4 value */
3138 assert(op == nir_intrinsic_image_deref_store);
3139 intrin->num_components = 4;
3140 intrin->src[3] = nir_src_for_ssa(expand_to_vec4(&b->nb, value));
3141 /* Only OpImageWrite can support a lod parameter if
3142 * SPV_AMD_shader_image_load_store_lod is used but the current NIR
3143 * intrinsics definition for atomics requires us to set it for
3144 * OpAtomicStore.
3145 */
3146 intrin->src[4] = nir_src_for_ssa(image.lod);
3147 break;
3148 }
3149
3150 case SpvOpAtomicCompareExchange:
3151 case SpvOpAtomicCompareExchangeWeak:
3152 case SpvOpAtomicIIncrement:
3153 case SpvOpAtomicIDecrement:
3154 case SpvOpAtomicExchange:
3155 case SpvOpAtomicIAdd:
3156 case SpvOpAtomicISub:
3157 case SpvOpAtomicSMin:
3158 case SpvOpAtomicUMin:
3159 case SpvOpAtomicSMax:
3160 case SpvOpAtomicUMax:
3161 case SpvOpAtomicAnd:
3162 case SpvOpAtomicOr:
3163 case SpvOpAtomicXor:
3164 case SpvOpAtomicFAddEXT:
3165 fill_common_atomic_sources(b, opcode, w, &intrin->src[3]);
3166 break;
3167
3168 default:
3169 vtn_fail_with_opcode("Invalid image opcode", opcode);
3170 }
3171
3172 /* Image operations implicitly have the Image storage memory semantics. */
3173 semantics |= SpvMemorySemanticsImageMemoryMask;
3174
3175 SpvMemorySemanticsMask before_semantics;
3176 SpvMemorySemanticsMask after_semantics;
3177 vtn_split_barrier_semantics(b, semantics, &before_semantics, &after_semantics);
3178
3179 if (before_semantics)
3180 vtn_emit_memory_barrier(b, scope, before_semantics);
3181
3182 if (opcode != SpvOpImageWrite && opcode != SpvOpAtomicStore) {
3183 struct vtn_type *type = vtn_get_type(b, w[1]);
3184
3185 unsigned dest_components = glsl_get_vector_elements(type->type);
3186 if (nir_intrinsic_infos[op].dest_components == 0)
3187 intrin->num_components = dest_components;
3188
3189 nir_ssa_dest_init(&intrin->instr, &intrin->dest,
3190 nir_intrinsic_dest_components(intrin), 32, NULL);
3191
3192 nir_builder_instr_insert(&b->nb, &intrin->instr);
3193
3194 nir_ssa_def *result = &intrin->dest.ssa;
3195 if (nir_intrinsic_dest_components(intrin) != dest_components)
3196 result = nir_channels(&b->nb, result, (1 << dest_components) - 1);
3197
3198 vtn_push_nir_ssa(b, w[2], result);
3199 } else {
3200 nir_builder_instr_insert(&b->nb, &intrin->instr);
3201 }
3202
3203 if (after_semantics)
3204 vtn_emit_memory_barrier(b, scope, after_semantics);
3205 }
3206
3207 static nir_intrinsic_op
3208 get_ssbo_nir_atomic_op(struct vtn_builder *b, SpvOp opcode)
3209 {
3210 switch (opcode) {
3211 case SpvOpAtomicLoad: return nir_intrinsic_load_ssbo;
3212 case SpvOpAtomicStore: return nir_intrinsic_store_ssbo;
3213 #define OP(S, N) case SpvOp##S: return nir_intrinsic_ssbo_##N;
3214 OP(AtomicExchange, atomic_exchange)
3215 OP(AtomicCompareExchange, atomic_comp_swap)
3216 OP(AtomicCompareExchangeWeak, atomic_comp_swap)
3217 OP(AtomicIIncrement, atomic_add)
3218 OP(AtomicIDecrement, atomic_add)
3219 OP(AtomicIAdd, atomic_add)
3220 OP(AtomicISub, atomic_add)
3221 OP(AtomicSMin, atomic_imin)
3222 OP(AtomicUMin, atomic_umin)
3223 OP(AtomicSMax, atomic_imax)
3224 OP(AtomicUMax, atomic_umax)
3225 OP(AtomicAnd, atomic_and)
3226 OP(AtomicOr, atomic_or)
3227 OP(AtomicXor, atomic_xor)
3228 OP(AtomicFAddEXT, atomic_fadd)
3229 #undef OP
3230 default:
3231 vtn_fail_with_opcode("Invalid SSBO atomic", opcode);
3232 }
3233 }
3234
3235 static nir_intrinsic_op
3236 get_uniform_nir_atomic_op(struct vtn_builder *b, SpvOp opcode)
3237 {
3238 switch (opcode) {
3239 #define OP(S, N) case SpvOp##S: return nir_intrinsic_atomic_counter_ ##N;
3240 OP(AtomicLoad, read_deref)
3241 OP(AtomicExchange, exchange)
3242 OP(AtomicCompareExchange, comp_swap)
3243 OP(AtomicCompareExchangeWeak, comp_swap)
3244 OP(AtomicIIncrement, inc_deref)
3245 OP(AtomicIDecrement, post_dec_deref)
3246 OP(AtomicIAdd, add_deref)
3247 OP(AtomicISub, add_deref)
3248 OP(AtomicUMin, min_deref)
3249 OP(AtomicUMax, max_deref)
3250 OP(AtomicAnd, and_deref)
3251 OP(AtomicOr, or_deref)
3252 OP(AtomicXor, xor_deref)
3253 #undef OP
3254 default:
3255 /* We left the following out: AtomicStore, AtomicSMin and
3256 * AtomicSmax. Right now there are not nir intrinsics for them. At this
3257 * moment Atomic Counter support is needed for ARB_spirv support, so is
3258 * only need to support GLSL Atomic Counters that are uints and don't
3259 * allow direct storage.
3260 */
3261 vtn_fail("Invalid uniform atomic");
3262 }
3263 }
3264
3265 static nir_intrinsic_op
3266 get_deref_nir_atomic_op(struct vtn_builder *b, SpvOp opcode)
3267 {
3268 switch (opcode) {
3269 case SpvOpAtomicLoad: return nir_intrinsic_load_deref;
3270 case SpvOpAtomicStore: return nir_intrinsic_store_deref;
3271 #define OP(S, N) case SpvOp##S: return nir_intrinsic_deref_##N;
3272 OP(AtomicExchange, atomic_exchange)
3273 OP(AtomicCompareExchange, atomic_comp_swap)
3274 OP(AtomicCompareExchangeWeak, atomic_comp_swap)
3275 OP(AtomicIIncrement, atomic_add)
3276 OP(AtomicIDecrement, atomic_add)
3277 OP(AtomicIAdd, atomic_add)
3278 OP(AtomicISub, atomic_add)
3279 OP(AtomicSMin, atomic_imin)
3280 OP(AtomicUMin, atomic_umin)
3281 OP(AtomicSMax, atomic_imax)
3282 OP(AtomicUMax, atomic_umax)
3283 OP(AtomicAnd, atomic_and)
3284 OP(AtomicOr, atomic_or)
3285 OP(AtomicXor, atomic_xor)
3286 OP(AtomicFAddEXT, atomic_fadd)
3287 #undef OP
3288 default:
3289 vtn_fail_with_opcode("Invalid shared atomic", opcode);
3290 }
3291 }
3292
3293 /*
3294 * Handles shared atomics, ssbo atomics and atomic counters.
3295 */
3296 static void
3297 vtn_handle_atomics(struct vtn_builder *b, SpvOp opcode,
3298 const uint32_t *w, UNUSED unsigned count)
3299 {
3300 struct vtn_pointer *ptr;
3301 nir_intrinsic_instr *atomic;
3302
3303 SpvScope scope = SpvScopeInvocation;
3304 SpvMemorySemanticsMask semantics = 0;
3305
3306 switch (opcode) {
3307 case SpvOpAtomicLoad:
3308 case SpvOpAtomicExchange:
3309 case SpvOpAtomicCompareExchange:
3310 case SpvOpAtomicCompareExchangeWeak:
3311 case SpvOpAtomicIIncrement:
3312 case SpvOpAtomicIDecrement:
3313 case SpvOpAtomicIAdd:
3314 case SpvOpAtomicISub:
3315 case SpvOpAtomicSMin:
3316 case SpvOpAtomicUMin:
3317 case SpvOpAtomicSMax:
3318 case SpvOpAtomicUMax:
3319 case SpvOpAtomicAnd:
3320 case SpvOpAtomicOr:
3321 case SpvOpAtomicXor:
3322 case SpvOpAtomicFAddEXT:
3323 ptr = vtn_value(b, w[3], vtn_value_type_pointer)->pointer;
3324 scope = vtn_constant_uint(b, w[4]);
3325 semantics = vtn_constant_uint(b, w[5]);
3326 break;
3327
3328 case SpvOpAtomicStore:
3329 ptr = vtn_value(b, w[1], vtn_value_type_pointer)->pointer;
3330 scope = vtn_constant_uint(b, w[2]);
3331 semantics = vtn_constant_uint(b, w[3]);
3332 break;
3333
3334 default:
3335 vtn_fail_with_opcode("Invalid SPIR-V atomic", opcode);
3336 }
3337
3338 /* uniform as "atomic counter uniform" */
3339 if (ptr->mode == vtn_variable_mode_atomic_counter) {
3340 nir_deref_instr *deref = vtn_pointer_to_deref(b, ptr);
3341 nir_intrinsic_op op = get_uniform_nir_atomic_op(b, opcode);
3342 atomic = nir_intrinsic_instr_create(b->nb.shader, op);
3343 atomic->src[0] = nir_src_for_ssa(&deref->dest.ssa);
3344
3345 /* SSBO needs to initialize index/offset. In this case we don't need to,
3346 * as that info is already stored on the ptr->var->var nir_variable (see
3347 * vtn_create_variable)
3348 */
3349
3350 switch (opcode) {
3351 case SpvOpAtomicLoad:
3352 case SpvOpAtomicExchange:
3353 case SpvOpAtomicCompareExchange:
3354 case SpvOpAtomicCompareExchangeWeak:
3355 case SpvOpAtomicIIncrement:
3356 case SpvOpAtomicIDecrement:
3357 case SpvOpAtomicIAdd:
3358 case SpvOpAtomicISub:
3359 case SpvOpAtomicSMin:
3360 case SpvOpAtomicUMin:
3361 case SpvOpAtomicSMax:
3362 case SpvOpAtomicUMax:
3363 case SpvOpAtomicAnd:
3364 case SpvOpAtomicOr:
3365 case SpvOpAtomicXor:
3366 /* Nothing: we don't need to call fill_common_atomic_sources here, as
3367 * atomic counter uniforms doesn't have sources
3368 */
3369 break;
3370
3371 default:
3372 unreachable("Invalid SPIR-V atomic");
3373
3374 }
3375 } else if (vtn_pointer_uses_ssa_offset(b, ptr)) {
3376 nir_ssa_def *offset, *index;
3377 offset = vtn_pointer_to_offset(b, ptr, &index);
3378
3379 assert(ptr->mode == vtn_variable_mode_ssbo);
3380
3381 nir_intrinsic_op op = get_ssbo_nir_atomic_op(b, opcode);
3382 atomic = nir_intrinsic_instr_create(b->nb.shader, op);
3383
3384 nir_intrinsic_set_access(atomic, ACCESS_COHERENT);
3385
3386 int src = 0;
3387 switch (opcode) {
3388 case SpvOpAtomicLoad:
3389 atomic->num_components = glsl_get_vector_elements(ptr->type->type);
3390 nir_intrinsic_set_align(atomic, 4, 0);
3391 if (ptr->mode == vtn_variable_mode_ssbo)
3392 atomic->src[src++] = nir_src_for_ssa(index);
3393 atomic->src[src++] = nir_src_for_ssa(offset);
3394 break;
3395
3396 case SpvOpAtomicStore:
3397 atomic->num_components = glsl_get_vector_elements(ptr->type->type);
3398 nir_intrinsic_set_write_mask(atomic, (1 << atomic->num_components) - 1);
3399 nir_intrinsic_set_align(atomic, 4, 0);
3400 atomic->src[src++] = nir_src_for_ssa(vtn_get_nir_ssa(b, w[4]));
3401 if (ptr->mode == vtn_variable_mode_ssbo)
3402 atomic->src[src++] = nir_src_for_ssa(index);
3403 atomic->src[src++] = nir_src_for_ssa(offset);
3404 break;
3405
3406 case SpvOpAtomicExchange:
3407 case SpvOpAtomicCompareExchange:
3408 case SpvOpAtomicCompareExchangeWeak:
3409 case SpvOpAtomicIIncrement:
3410 case SpvOpAtomicIDecrement:
3411 case SpvOpAtomicIAdd:
3412 case SpvOpAtomicISub:
3413 case SpvOpAtomicSMin:
3414 case SpvOpAtomicUMin:
3415 case SpvOpAtomicSMax:
3416 case SpvOpAtomicUMax:
3417 case SpvOpAtomicAnd:
3418 case SpvOpAtomicOr:
3419 case SpvOpAtomicXor:
3420 case SpvOpAtomicFAddEXT:
3421 if (ptr->mode == vtn_variable_mode_ssbo)
3422 atomic->src[src++] = nir_src_for_ssa(index);
3423 atomic->src[src++] = nir_src_for_ssa(offset);
3424 fill_common_atomic_sources(b, opcode, w, &atomic->src[src]);
3425 break;
3426
3427 default:
3428 vtn_fail_with_opcode("Invalid SPIR-V atomic", opcode);
3429 }
3430 } else {
3431 nir_deref_instr *deref = vtn_pointer_to_deref(b, ptr);
3432 const struct glsl_type *deref_type = deref->type;
3433 nir_intrinsic_op op = get_deref_nir_atomic_op(b, opcode);
3434 atomic = nir_intrinsic_instr_create(b->nb.shader, op);
3435 atomic->src[0] = nir_src_for_ssa(&deref->dest.ssa);
3436
3437 if (ptr->mode != vtn_variable_mode_workgroup)
3438 nir_intrinsic_set_access(atomic, ACCESS_COHERENT);
3439
3440 switch (opcode) {
3441 case SpvOpAtomicLoad:
3442 atomic->num_components = glsl_get_vector_elements(deref_type);
3443 break;
3444
3445 case SpvOpAtomicStore:
3446 atomic->num_components = glsl_get_vector_elements(deref_type);
3447 nir_intrinsic_set_write_mask(atomic, (1 << atomic->num_components) - 1);
3448 atomic->src[1] = nir_src_for_ssa(vtn_get_nir_ssa(b, w[4]));
3449 break;
3450
3451 case SpvOpAtomicExchange:
3452 case SpvOpAtomicCompareExchange:
3453 case SpvOpAtomicCompareExchangeWeak:
3454 case SpvOpAtomicIIncrement:
3455 case SpvOpAtomicIDecrement:
3456 case SpvOpAtomicIAdd:
3457 case SpvOpAtomicISub:
3458 case SpvOpAtomicSMin:
3459 case SpvOpAtomicUMin:
3460 case SpvOpAtomicSMax:
3461 case SpvOpAtomicUMax:
3462 case SpvOpAtomicAnd:
3463 case SpvOpAtomicOr:
3464 case SpvOpAtomicXor:
3465 case SpvOpAtomicFAddEXT:
3466 fill_common_atomic_sources(b, opcode, w, &atomic->src[1]);
3467 break;
3468
3469 default:
3470 vtn_fail_with_opcode("Invalid SPIR-V atomic", opcode);
3471 }
3472 }
3473
3474 /* Atomic ordering operations will implicitly apply to the atomic operation
3475 * storage class, so include that too.
3476 */
3477 semantics |= vtn_storage_class_to_memory_semantics(ptr->ptr_type->storage_class);
3478
3479 SpvMemorySemanticsMask before_semantics;
3480 SpvMemorySemanticsMask after_semantics;
3481 vtn_split_barrier_semantics(b, semantics, &before_semantics, &after_semantics);
3482
3483 if (before_semantics)
3484 vtn_emit_memory_barrier(b, scope, before_semantics);
3485
3486 if (opcode != SpvOpAtomicStore) {
3487 struct vtn_type *type = vtn_get_type(b, w[1]);
3488
3489 nir_ssa_dest_init(&atomic->instr, &atomic->dest,
3490 glsl_get_vector_elements(type->type),
3491 glsl_get_bit_size(type->type), NULL);
3492
3493 vtn_push_nir_ssa(b, w[2], &atomic->dest.ssa);
3494 }
3495
3496 nir_builder_instr_insert(&b->nb, &atomic->instr);
3497
3498 if (after_semantics)
3499 vtn_emit_memory_barrier(b, scope, after_semantics);
3500 }
3501
3502 static nir_alu_instr *
3503 create_vec(struct vtn_builder *b, unsigned num_components, unsigned bit_size)
3504 {
3505 nir_op op = nir_op_vec(num_components);
3506 nir_alu_instr *vec = nir_alu_instr_create(b->shader, op);
3507 nir_ssa_dest_init(&vec->instr, &vec->dest.dest, num_components,
3508 bit_size, NULL);
3509 vec->dest.write_mask = (1 << num_components) - 1;
3510
3511 return vec;
3512 }
3513
3514 struct vtn_ssa_value *
3515 vtn_ssa_transpose(struct vtn_builder *b, struct vtn_ssa_value *src)
3516 {
3517 if (src->transposed)
3518 return src->transposed;
3519
3520 struct vtn_ssa_value *dest =
3521 vtn_create_ssa_value(b, glsl_transposed_type(src->type));
3522
3523 for (unsigned i = 0; i < glsl_get_matrix_columns(dest->type); i++) {
3524 nir_alu_instr *vec = create_vec(b, glsl_get_matrix_columns(src->type),
3525 glsl_get_bit_size(src->type));
3526 if (glsl_type_is_vector_or_scalar(src->type)) {
3527 vec->src[0].src = nir_src_for_ssa(src->def);
3528 vec->src[0].swizzle[0] = i;
3529 } else {
3530 for (unsigned j = 0; j < glsl_get_matrix_columns(src->type); j++) {
3531 vec->src[j].src = nir_src_for_ssa(src->elems[j]->def);
3532 vec->src[j].swizzle[0] = i;
3533 }
3534 }
3535 nir_builder_instr_insert(&b->nb, &vec->instr);
3536 dest->elems[i]->def = &vec->dest.dest.ssa;
3537 }
3538
3539 dest->transposed = src;
3540
3541 return dest;
3542 }
3543
3544 static nir_ssa_def *
3545 vtn_vector_shuffle(struct vtn_builder *b, unsigned num_components,
3546 nir_ssa_def *src0, nir_ssa_def *src1,
3547 const uint32_t *indices)
3548 {
3549 nir_alu_instr *vec = create_vec(b, num_components, src0->bit_size);
3550
3551 for (unsigned i = 0; i < num_components; i++) {
3552 uint32_t index = indices[i];
3553 if (index == 0xffffffff) {
3554 vec->src[i].src =
3555 nir_src_for_ssa(nir_ssa_undef(&b->nb, 1, src0->bit_size));
3556 } else if (index < src0->num_components) {
3557 vec->src[i].src = nir_src_for_ssa(src0);
3558 vec->src[i].swizzle[0] = index;
3559 } else {
3560 vec->src[i].src = nir_src_for_ssa(src1);
3561 vec->src[i].swizzle[0] = index - src0->num_components;
3562 }
3563 }
3564
3565 nir_builder_instr_insert(&b->nb, &vec->instr);
3566
3567 return &vec->dest.dest.ssa;
3568 }
3569
3570 /*
3571 * Concatentates a number of vectors/scalars together to produce a vector