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