spirv: Drop the OpenCL type layout code
[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 case SpvDecorationCPacked:
1045 vtn_warn("Decoration not allowed on struct members: %s",
1046 spirv_decoration_to_string(dec->decoration));
1047 break;
1048
1049 case SpvDecorationXfbBuffer:
1050 case SpvDecorationXfbStride:
1051 /* This is handled later by var_decoration_cb in vtn_variables.c */
1052 break;
1053
1054 case SpvDecorationSaturatedConversion:
1055 case SpvDecorationFuncParamAttr:
1056 case SpvDecorationFPRoundingMode:
1057 case SpvDecorationFPFastMathMode:
1058 case SpvDecorationAlignment:
1059 if (b->shader->info.stage != MESA_SHADER_KERNEL) {
1060 vtn_warn("Decoration only allowed for CL-style kernels: %s",
1061 spirv_decoration_to_string(dec->decoration));
1062 }
1063 break;
1064
1065 case SpvDecorationUserSemantic:
1066 case SpvDecorationUserTypeGOOGLE:
1067 /* User semantic decorations can safely be ignored by the driver. */
1068 break;
1069
1070 default:
1071 vtn_fail_with_decoration("Unhandled decoration", dec->decoration);
1072 }
1073 }
1074
1075 /** Chases the array type all the way down to the tail and rewrites the
1076 * glsl_types to be based off the tail's glsl_type.
1077 */
1078 static void
1079 vtn_array_type_rewrite_glsl_type(struct vtn_type *type)
1080 {
1081 if (type->base_type != vtn_base_type_array)
1082 return;
1083
1084 vtn_array_type_rewrite_glsl_type(type->array_element);
1085
1086 type->type = glsl_array_type(type->array_element->type,
1087 type->length, type->stride);
1088 }
1089
1090 /* Matrix strides are handled as a separate pass because we need to know
1091 * whether the matrix is row-major or not first.
1092 */
1093 static void
1094 struct_member_matrix_stride_cb(struct vtn_builder *b,
1095 UNUSED struct vtn_value *val, int member,
1096 const struct vtn_decoration *dec,
1097 void *void_ctx)
1098 {
1099 if (dec->decoration != SpvDecorationMatrixStride)
1100 return;
1101
1102 vtn_fail_if(member < 0,
1103 "The MatrixStride decoration is only allowed on members "
1104 "of OpTypeStruct");
1105 vtn_fail_if(dec->operands[0] == 0, "MatrixStride must be non-zero");
1106
1107 struct member_decoration_ctx *ctx = void_ctx;
1108
1109 struct vtn_type *mat_type = mutable_matrix_member(b, ctx->type, member);
1110 if (mat_type->row_major) {
1111 mat_type->array_element = vtn_type_copy(b, mat_type->array_element);
1112 mat_type->stride = mat_type->array_element->stride;
1113 mat_type->array_element->stride = dec->operands[0];
1114
1115 mat_type->type = glsl_explicit_matrix_type(mat_type->type,
1116 dec->operands[0], true);
1117 mat_type->array_element->type = glsl_get_column_type(mat_type->type);
1118 } else {
1119 vtn_assert(mat_type->array_element->stride > 0);
1120 mat_type->stride = dec->operands[0];
1121
1122 mat_type->type = glsl_explicit_matrix_type(mat_type->type,
1123 dec->operands[0], false);
1124 }
1125
1126 /* Now that we've replaced the glsl_type with a properly strided matrix
1127 * type, rewrite the member type so that it's an array of the proper kind
1128 * of glsl_type.
1129 */
1130 vtn_array_type_rewrite_glsl_type(ctx->type->members[member]);
1131 ctx->fields[member].type = ctx->type->members[member]->type;
1132 }
1133
1134 static void
1135 struct_packed_decoration_cb(struct vtn_builder *b,
1136 struct vtn_value *val, int member,
1137 const struct vtn_decoration *dec, void *void_ctx)
1138 {
1139 vtn_assert(val->type->base_type == vtn_base_type_struct);
1140 if (dec->decoration == SpvDecorationCPacked) {
1141 if (b->shader->info.stage != MESA_SHADER_KERNEL) {
1142 vtn_warn("Decoration only allowed for CL-style kernels: %s",
1143 spirv_decoration_to_string(dec->decoration));
1144 }
1145 val->type->packed = true;
1146 }
1147 }
1148
1149 static void
1150 struct_block_decoration_cb(struct vtn_builder *b,
1151 struct vtn_value *val, int member,
1152 const struct vtn_decoration *dec, void *ctx)
1153 {
1154 if (member != -1)
1155 return;
1156
1157 struct vtn_type *type = val->type;
1158 if (dec->decoration == SpvDecorationBlock)
1159 type->block = true;
1160 else if (dec->decoration == SpvDecorationBufferBlock)
1161 type->buffer_block = true;
1162 }
1163
1164 static void
1165 type_decoration_cb(struct vtn_builder *b,
1166 struct vtn_value *val, int member,
1167 const struct vtn_decoration *dec, UNUSED void *ctx)
1168 {
1169 struct vtn_type *type = val->type;
1170
1171 if (member != -1) {
1172 /* This should have been handled by OpTypeStruct */
1173 assert(val->type->base_type == vtn_base_type_struct);
1174 assert(member >= 0 && member < val->type->length);
1175 return;
1176 }
1177
1178 switch (dec->decoration) {
1179 case SpvDecorationArrayStride:
1180 vtn_assert(type->base_type == vtn_base_type_array ||
1181 type->base_type == vtn_base_type_pointer);
1182 break;
1183 case SpvDecorationBlock:
1184 vtn_assert(type->base_type == vtn_base_type_struct);
1185 vtn_assert(type->block);
1186 break;
1187 case SpvDecorationBufferBlock:
1188 vtn_assert(type->base_type == vtn_base_type_struct);
1189 vtn_assert(type->buffer_block);
1190 break;
1191 case SpvDecorationGLSLShared:
1192 case SpvDecorationGLSLPacked:
1193 /* Ignore these, since we get explicit offsets anyways */
1194 break;
1195
1196 case SpvDecorationRowMajor:
1197 case SpvDecorationColMajor:
1198 case SpvDecorationMatrixStride:
1199 case SpvDecorationBuiltIn:
1200 case SpvDecorationNoPerspective:
1201 case SpvDecorationFlat:
1202 case SpvDecorationPatch:
1203 case SpvDecorationCentroid:
1204 case SpvDecorationSample:
1205 case SpvDecorationExplicitInterpAMD:
1206 case SpvDecorationVolatile:
1207 case SpvDecorationCoherent:
1208 case SpvDecorationNonWritable:
1209 case SpvDecorationNonReadable:
1210 case SpvDecorationUniform:
1211 case SpvDecorationUniformId:
1212 case SpvDecorationLocation:
1213 case SpvDecorationComponent:
1214 case SpvDecorationOffset:
1215 case SpvDecorationXfbBuffer:
1216 case SpvDecorationXfbStride:
1217 case SpvDecorationUserSemantic:
1218 vtn_warn("Decoration only allowed for struct members: %s",
1219 spirv_decoration_to_string(dec->decoration));
1220 break;
1221
1222 case SpvDecorationStream:
1223 /* We don't need to do anything here, as stream is filled up when
1224 * aplying the decoration to a variable, just check that if it is not a
1225 * struct member, it should be a struct.
1226 */
1227 vtn_assert(type->base_type == vtn_base_type_struct);
1228 break;
1229
1230 case SpvDecorationRelaxedPrecision:
1231 case SpvDecorationSpecId:
1232 case SpvDecorationInvariant:
1233 case SpvDecorationRestrict:
1234 case SpvDecorationAliased:
1235 case SpvDecorationConstant:
1236 case SpvDecorationIndex:
1237 case SpvDecorationBinding:
1238 case SpvDecorationDescriptorSet:
1239 case SpvDecorationLinkageAttributes:
1240 case SpvDecorationNoContraction:
1241 case SpvDecorationInputAttachmentIndex:
1242 vtn_warn("Decoration not allowed on types: %s",
1243 spirv_decoration_to_string(dec->decoration));
1244 break;
1245
1246 case SpvDecorationCPacked:
1247 /* Handled when parsing a struct type, nothing to do here. */
1248 break;
1249
1250 case SpvDecorationSaturatedConversion:
1251 case SpvDecorationFuncParamAttr:
1252 case SpvDecorationFPRoundingMode:
1253 case SpvDecorationFPFastMathMode:
1254 case SpvDecorationAlignment:
1255 vtn_warn("Decoration only allowed for CL-style kernels: %s",
1256 spirv_decoration_to_string(dec->decoration));
1257 break;
1258
1259 case SpvDecorationUserTypeGOOGLE:
1260 /* User semantic decorations can safely be ignored by the driver. */
1261 break;
1262
1263 default:
1264 vtn_fail_with_decoration("Unhandled decoration", dec->decoration);
1265 }
1266 }
1267
1268 static unsigned
1269 translate_image_format(struct vtn_builder *b, SpvImageFormat format)
1270 {
1271 switch (format) {
1272 case SpvImageFormatUnknown: return PIPE_FORMAT_NONE;
1273 case SpvImageFormatRgba32f: return PIPE_FORMAT_R32G32B32A32_FLOAT;
1274 case SpvImageFormatRgba16f: return PIPE_FORMAT_R16G16B16A16_FLOAT;
1275 case SpvImageFormatR32f: return PIPE_FORMAT_R32_FLOAT;
1276 case SpvImageFormatRgba8: return PIPE_FORMAT_R8G8B8A8_UNORM;
1277 case SpvImageFormatRgba8Snorm: return PIPE_FORMAT_R8G8B8A8_SNORM;
1278 case SpvImageFormatRg32f: return PIPE_FORMAT_R32G32_FLOAT;
1279 case SpvImageFormatRg16f: return PIPE_FORMAT_R16G16_FLOAT;
1280 case SpvImageFormatR11fG11fB10f: return PIPE_FORMAT_R11G11B10_FLOAT;
1281 case SpvImageFormatR16f: return PIPE_FORMAT_R16_FLOAT;
1282 case SpvImageFormatRgba16: return PIPE_FORMAT_R16G16B16A16_UNORM;
1283 case SpvImageFormatRgb10A2: return PIPE_FORMAT_R10G10B10A2_UNORM;
1284 case SpvImageFormatRg16: return PIPE_FORMAT_R16G16_UNORM;
1285 case SpvImageFormatRg8: return PIPE_FORMAT_R8G8_UNORM;
1286 case SpvImageFormatR16: return PIPE_FORMAT_R16_UNORM;
1287 case SpvImageFormatR8: return PIPE_FORMAT_R8_UNORM;
1288 case SpvImageFormatRgba16Snorm: return PIPE_FORMAT_R16G16B16A16_SNORM;
1289 case SpvImageFormatRg16Snorm: return PIPE_FORMAT_R16G16_SNORM;
1290 case SpvImageFormatRg8Snorm: return PIPE_FORMAT_R8G8_SNORM;
1291 case SpvImageFormatR16Snorm: return PIPE_FORMAT_R16_SNORM;
1292 case SpvImageFormatR8Snorm: return PIPE_FORMAT_R8_SNORM;
1293 case SpvImageFormatRgba32i: return PIPE_FORMAT_R32G32B32A32_SINT;
1294 case SpvImageFormatRgba16i: return PIPE_FORMAT_R16G16B16A16_SINT;
1295 case SpvImageFormatRgba8i: return PIPE_FORMAT_R8G8B8A8_SINT;
1296 case SpvImageFormatR32i: return PIPE_FORMAT_R32_SINT;
1297 case SpvImageFormatRg32i: return PIPE_FORMAT_R32G32_SINT;
1298 case SpvImageFormatRg16i: return PIPE_FORMAT_R16G16_SINT;
1299 case SpvImageFormatRg8i: return PIPE_FORMAT_R8G8_SINT;
1300 case SpvImageFormatR16i: return PIPE_FORMAT_R16_SINT;
1301 case SpvImageFormatR8i: return PIPE_FORMAT_R8_SINT;
1302 case SpvImageFormatRgba32ui: return PIPE_FORMAT_R32G32B32A32_UINT;
1303 case SpvImageFormatRgba16ui: return PIPE_FORMAT_R16G16B16A16_UINT;
1304 case SpvImageFormatRgba8ui: return PIPE_FORMAT_R8G8B8A8_UINT;
1305 case SpvImageFormatR32ui: return PIPE_FORMAT_R32_UINT;
1306 case SpvImageFormatRgb10a2ui: return PIPE_FORMAT_R10G10B10A2_UINT;
1307 case SpvImageFormatRg32ui: return PIPE_FORMAT_R32G32_UINT;
1308 case SpvImageFormatRg16ui: return PIPE_FORMAT_R16G16_UINT;
1309 case SpvImageFormatRg8ui: return PIPE_FORMAT_R8G8_UINT;
1310 case SpvImageFormatR16ui: return PIPE_FORMAT_R16_UINT;
1311 case SpvImageFormatR8ui: return PIPE_FORMAT_R8_UINT;
1312 default:
1313 vtn_fail("Invalid image format: %s (%u)",
1314 spirv_imageformat_to_string(format), format);
1315 }
1316 }
1317
1318 static void
1319 vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
1320 const uint32_t *w, unsigned count)
1321 {
1322 struct vtn_value *val = NULL;
1323
1324 /* In order to properly handle forward declarations, we have to defer
1325 * allocation for pointer types.
1326 */
1327 if (opcode != SpvOpTypePointer && opcode != SpvOpTypeForwardPointer) {
1328 val = vtn_push_value(b, w[1], vtn_value_type_type);
1329 vtn_fail_if(val->type != NULL,
1330 "Only pointers can have forward declarations");
1331 val->type = rzalloc(b, struct vtn_type);
1332 val->type->id = w[1];
1333 }
1334
1335 switch (opcode) {
1336 case SpvOpTypeVoid:
1337 val->type->base_type = vtn_base_type_void;
1338 val->type->type = glsl_void_type();
1339 break;
1340 case SpvOpTypeBool:
1341 val->type->base_type = vtn_base_type_scalar;
1342 val->type->type = glsl_bool_type();
1343 val->type->length = 1;
1344 break;
1345 case SpvOpTypeInt: {
1346 int bit_size = w[2];
1347 const bool signedness = w[3];
1348 val->type->base_type = vtn_base_type_scalar;
1349 switch (bit_size) {
1350 case 64:
1351 val->type->type = (signedness ? glsl_int64_t_type() : glsl_uint64_t_type());
1352 break;
1353 case 32:
1354 val->type->type = (signedness ? glsl_int_type() : glsl_uint_type());
1355 break;
1356 case 16:
1357 val->type->type = (signedness ? glsl_int16_t_type() : glsl_uint16_t_type());
1358 break;
1359 case 8:
1360 val->type->type = (signedness ? glsl_int8_t_type() : glsl_uint8_t_type());
1361 break;
1362 default:
1363 vtn_fail("Invalid int bit size: %u", bit_size);
1364 }
1365 val->type->length = 1;
1366 break;
1367 }
1368
1369 case SpvOpTypeFloat: {
1370 int bit_size = w[2];
1371 val->type->base_type = vtn_base_type_scalar;
1372 switch (bit_size) {
1373 case 16:
1374 val->type->type = glsl_float16_t_type();
1375 break;
1376 case 32:
1377 val->type->type = glsl_float_type();
1378 break;
1379 case 64:
1380 val->type->type = glsl_double_type();
1381 break;
1382 default:
1383 vtn_fail("Invalid float bit size: %u", bit_size);
1384 }
1385 val->type->length = 1;
1386 break;
1387 }
1388
1389 case SpvOpTypeVector: {
1390 struct vtn_type *base = vtn_get_type(b, w[2]);
1391 unsigned elems = w[3];
1392
1393 vtn_fail_if(base->base_type != vtn_base_type_scalar,
1394 "Base type for OpTypeVector must be a scalar");
1395 vtn_fail_if((elems < 2 || elems > 4) && (elems != 8) && (elems != 16),
1396 "Invalid component count for OpTypeVector");
1397
1398 val->type->base_type = vtn_base_type_vector;
1399 val->type->type = glsl_vector_type(glsl_get_base_type(base->type), elems);
1400 val->type->length = elems;
1401 val->type->stride = glsl_type_is_boolean(val->type->type)
1402 ? 4 : glsl_get_bit_size(base->type) / 8;
1403 val->type->array_element = base;
1404 break;
1405 }
1406
1407 case SpvOpTypeMatrix: {
1408 struct vtn_type *base = vtn_get_type(b, w[2]);
1409 unsigned columns = w[3];
1410
1411 vtn_fail_if(base->base_type != vtn_base_type_vector,
1412 "Base type for OpTypeMatrix must be a vector");
1413 vtn_fail_if(columns < 2 || columns > 4,
1414 "Invalid column count for OpTypeMatrix");
1415
1416 val->type->base_type = vtn_base_type_matrix;
1417 val->type->type = glsl_matrix_type(glsl_get_base_type(base->type),
1418 glsl_get_vector_elements(base->type),
1419 columns);
1420 vtn_fail_if(glsl_type_is_error(val->type->type),
1421 "Unsupported base type for OpTypeMatrix");
1422 assert(!glsl_type_is_error(val->type->type));
1423 val->type->length = columns;
1424 val->type->array_element = base;
1425 val->type->row_major = false;
1426 val->type->stride = 0;
1427 break;
1428 }
1429
1430 case SpvOpTypeRuntimeArray:
1431 case SpvOpTypeArray: {
1432 struct vtn_type *array_element = vtn_get_type(b, w[2]);
1433
1434 if (opcode == SpvOpTypeRuntimeArray) {
1435 /* A length of 0 is used to denote unsized arrays */
1436 val->type->length = 0;
1437 } else {
1438 val->type->length = vtn_constant_uint(b, w[3]);
1439 }
1440
1441 val->type->base_type = vtn_base_type_array;
1442 val->type->array_element = array_element;
1443
1444 vtn_foreach_decoration(b, val, array_stride_decoration_cb, NULL);
1445 val->type->type = glsl_array_type(array_element->type, val->type->length,
1446 val->type->stride);
1447 break;
1448 }
1449
1450 case SpvOpTypeStruct: {
1451 unsigned num_fields = count - 2;
1452 val->type->base_type = vtn_base_type_struct;
1453 val->type->length = num_fields;
1454 val->type->members = ralloc_array(b, struct vtn_type *, num_fields);
1455 val->type->offsets = ralloc_array(b, unsigned, num_fields);
1456 val->type->packed = false;
1457
1458 NIR_VLA(struct glsl_struct_field, fields, count);
1459 for (unsigned i = 0; i < num_fields; i++) {
1460 val->type->members[i] = vtn_get_type(b, w[i + 2]);
1461 fields[i] = (struct glsl_struct_field) {
1462 .type = val->type->members[i]->type,
1463 .name = ralloc_asprintf(b, "field%d", i),
1464 .location = -1,
1465 .offset = -1,
1466 };
1467 }
1468
1469 vtn_foreach_decoration(b, val, struct_packed_decoration_cb, NULL);
1470
1471 struct member_decoration_ctx ctx = {
1472 .num_fields = num_fields,
1473 .fields = fields,
1474 .type = val->type
1475 };
1476
1477 vtn_foreach_decoration(b, val, struct_member_decoration_cb, &ctx);
1478 vtn_foreach_decoration(b, val, struct_member_matrix_stride_cb, &ctx);
1479
1480 vtn_foreach_decoration(b, val, struct_block_decoration_cb, NULL);
1481
1482 const char *name = val->name;
1483
1484 if (val->type->block || val->type->buffer_block) {
1485 /* Packing will be ignored since types coming from SPIR-V are
1486 * explicitly laid out.
1487 */
1488 val->type->type = glsl_interface_type(fields, num_fields,
1489 /* packing */ 0, false,
1490 name ? name : "block");
1491 } else {
1492 val->type->type = glsl_struct_type(fields, num_fields,
1493 name ? name : "struct",
1494 val->type->packed);
1495 }
1496 break;
1497 }
1498
1499 case SpvOpTypeFunction: {
1500 val->type->base_type = vtn_base_type_function;
1501 val->type->type = NULL;
1502
1503 val->type->return_type = vtn_get_type(b, w[2]);
1504
1505 const unsigned num_params = count - 3;
1506 val->type->length = num_params;
1507 val->type->params = ralloc_array(b, struct vtn_type *, num_params);
1508 for (unsigned i = 0; i < count - 3; i++) {
1509 val->type->params[i] = vtn_get_type(b, w[i + 3]);
1510 }
1511 break;
1512 }
1513
1514 case SpvOpTypePointer:
1515 case SpvOpTypeForwardPointer: {
1516 /* We can't blindly push the value because it might be a forward
1517 * declaration.
1518 */
1519 val = vtn_untyped_value(b, w[1]);
1520
1521 SpvStorageClass storage_class = w[2];
1522
1523 if (val->value_type == vtn_value_type_invalid) {
1524 val->value_type = vtn_value_type_type;
1525 val->type = rzalloc(b, struct vtn_type);
1526 val->type->id = w[1];
1527 val->type->base_type = vtn_base_type_pointer;
1528 val->type->storage_class = storage_class;
1529
1530 /* These can actually be stored to nir_variables and used as SSA
1531 * values so they need a real glsl_type.
1532 */
1533 enum vtn_variable_mode mode = vtn_storage_class_to_mode(
1534 b, storage_class, NULL, NULL);
1535 val->type->type = nir_address_format_to_glsl_type(
1536 vtn_mode_to_address_format(b, mode));
1537 } else {
1538 vtn_fail_if(val->type->storage_class != storage_class,
1539 "The storage classes of an OpTypePointer and any "
1540 "OpTypeForwardPointers that provide forward "
1541 "declarations of it must match.");
1542 }
1543
1544 if (opcode == SpvOpTypePointer) {
1545 vtn_fail_if(val->type->deref != NULL,
1546 "While OpTypeForwardPointer can be used to provide a "
1547 "forward declaration of a pointer, OpTypePointer can "
1548 "only be used once for a given id.");
1549
1550 val->type->deref = vtn_get_type(b, w[3]);
1551
1552 /* Only certain storage classes use ArrayStride. The others (in
1553 * particular Workgroup) are expected to be laid out by the driver.
1554 */
1555 switch (storage_class) {
1556 case SpvStorageClassUniform:
1557 case SpvStorageClassPushConstant:
1558 case SpvStorageClassStorageBuffer:
1559 case SpvStorageClassPhysicalStorageBuffer:
1560 vtn_foreach_decoration(b, val, array_stride_decoration_cb, NULL);
1561 break;
1562 default:
1563 /* Nothing to do. */
1564 break;
1565 }
1566 }
1567 break;
1568 }
1569
1570 case SpvOpTypeImage: {
1571 val->type->base_type = vtn_base_type_image;
1572
1573 /* Images are represented in NIR as a scalar SSA value that is the
1574 * result of a deref instruction. An OpLoad on an OpTypeImage pointer
1575 * from UniformConstant memory just takes the NIR deref from the pointer
1576 * and turns it into an SSA value.
1577 */
1578 val->type->type = nir_address_format_to_glsl_type(
1579 vtn_mode_to_address_format(b, vtn_variable_mode_function));
1580
1581 const struct vtn_type *sampled_type = vtn_get_type(b, w[2]);
1582 if (b->shader->info.stage == MESA_SHADER_KERNEL) {
1583 vtn_fail_if(sampled_type->base_type != vtn_base_type_void,
1584 "Sampled type of OpTypeImage must be void for kernels");
1585 } else {
1586 vtn_fail_if(sampled_type->base_type != vtn_base_type_scalar ||
1587 glsl_get_bit_size(sampled_type->type) != 32,
1588 "Sampled type of OpTypeImage must be a 32-bit scalar");
1589 }
1590
1591 enum glsl_sampler_dim dim;
1592 switch ((SpvDim)w[3]) {
1593 case SpvDim1D: dim = GLSL_SAMPLER_DIM_1D; break;
1594 case SpvDim2D: dim = GLSL_SAMPLER_DIM_2D; break;
1595 case SpvDim3D: dim = GLSL_SAMPLER_DIM_3D; break;
1596 case SpvDimCube: dim = GLSL_SAMPLER_DIM_CUBE; break;
1597 case SpvDimRect: dim = GLSL_SAMPLER_DIM_RECT; break;
1598 case SpvDimBuffer: dim = GLSL_SAMPLER_DIM_BUF; break;
1599 case SpvDimSubpassData: dim = GLSL_SAMPLER_DIM_SUBPASS; break;
1600 default:
1601 vtn_fail("Invalid SPIR-V image dimensionality: %s (%u)",
1602 spirv_dim_to_string((SpvDim)w[3]), w[3]);
1603 }
1604
1605 /* w[4]: as per Vulkan spec "Validation Rules within a Module",
1606 * The “Depth” operand of OpTypeImage is ignored.
1607 */
1608 bool is_array = w[5];
1609 bool multisampled = w[6];
1610 unsigned sampled = w[7];
1611 SpvImageFormat format = w[8];
1612
1613 if (count > 9)
1614 val->type->access_qualifier = w[9];
1615 else if (b->shader->info.stage == MESA_SHADER_KERNEL)
1616 /* Per the CL C spec: If no qualifier is provided, read_only is assumed. */
1617 val->type->access_qualifier = SpvAccessQualifierReadOnly;
1618 else
1619 val->type->access_qualifier = SpvAccessQualifierReadWrite;
1620
1621 if (multisampled) {
1622 if (dim == GLSL_SAMPLER_DIM_2D)
1623 dim = GLSL_SAMPLER_DIM_MS;
1624 else if (dim == GLSL_SAMPLER_DIM_SUBPASS)
1625 dim = GLSL_SAMPLER_DIM_SUBPASS_MS;
1626 else
1627 vtn_fail("Unsupported multisampled image type");
1628 }
1629
1630 val->type->image_format = translate_image_format(b, format);
1631
1632 enum glsl_base_type sampled_base_type =
1633 glsl_get_base_type(sampled_type->type);
1634 if (sampled == 1) {
1635 val->type->glsl_image = glsl_sampler_type(dim, false, is_array,
1636 sampled_base_type);
1637 } else if (sampled == 2) {
1638 val->type->glsl_image = glsl_image_type(dim, is_array,
1639 sampled_base_type);
1640 } else if (b->shader->info.stage == MESA_SHADER_KERNEL) {
1641 val->type->glsl_image = glsl_image_type(dim, is_array,
1642 GLSL_TYPE_VOID);
1643 } else {
1644 vtn_fail("We need to know if the image will be sampled");
1645 }
1646 break;
1647 }
1648
1649 case SpvOpTypeSampledImage: {
1650 val->type->base_type = vtn_base_type_sampled_image;
1651 val->type->image = vtn_get_type(b, w[2]);
1652
1653 /* Sampled images are represented NIR as a vec2 SSA value where each
1654 * component is the result of a deref instruction. The first component
1655 * is the image and the second is the sampler. An OpLoad on an
1656 * OpTypeSampledImage pointer from UniformConstant memory just takes
1657 * the NIR deref from the pointer and duplicates it to both vector
1658 * components.
1659 */
1660 nir_address_format addr_format =
1661 vtn_mode_to_address_format(b, vtn_variable_mode_function);
1662 assert(nir_address_format_num_components(addr_format) == 1);
1663 unsigned bit_size = nir_address_format_bit_size(addr_format);
1664 assert(bit_size == 32 || bit_size == 64);
1665
1666 enum glsl_base_type base_type =
1667 bit_size == 32 ? GLSL_TYPE_UINT : GLSL_TYPE_UINT64;
1668 val->type->type = glsl_vector_type(base_type, 2);
1669 break;
1670 }
1671
1672 case SpvOpTypeSampler:
1673 val->type->base_type = vtn_base_type_sampler;
1674
1675 /* Samplers are represented in NIR as a scalar SSA value that is the
1676 * result of a deref instruction. An OpLoad on an OpTypeSampler pointer
1677 * from UniformConstant memory just takes the NIR deref from the pointer
1678 * and turns it into an SSA value.
1679 */
1680 val->type->type = nir_address_format_to_glsl_type(
1681 vtn_mode_to_address_format(b, vtn_variable_mode_function));
1682 break;
1683
1684 case SpvOpTypeOpaque:
1685 case SpvOpTypeEvent:
1686 case SpvOpTypeDeviceEvent:
1687 case SpvOpTypeReserveId:
1688 case SpvOpTypeQueue:
1689 case SpvOpTypePipe:
1690 default:
1691 vtn_fail_with_opcode("Unhandled opcode", opcode);
1692 }
1693
1694 vtn_foreach_decoration(b, val, type_decoration_cb, NULL);
1695
1696 if (val->type->base_type == vtn_base_type_struct &&
1697 (val->type->block || val->type->buffer_block)) {
1698 for (unsigned i = 0; i < val->type->length; i++) {
1699 vtn_fail_if(vtn_type_contains_block(b, val->type->members[i]),
1700 "Block and BufferBlock decorations cannot decorate a "
1701 "structure type that is nested at any level inside "
1702 "another structure type decorated with Block or "
1703 "BufferBlock.");
1704 }
1705 }
1706 }
1707
1708 static nir_constant *
1709 vtn_null_constant(struct vtn_builder *b, struct vtn_type *type)
1710 {
1711 nir_constant *c = rzalloc(b, nir_constant);
1712
1713 switch (type->base_type) {
1714 case vtn_base_type_scalar:
1715 case vtn_base_type_vector:
1716 /* Nothing to do here. It's already initialized to zero */
1717 break;
1718
1719 case vtn_base_type_pointer: {
1720 enum vtn_variable_mode mode = vtn_storage_class_to_mode(
1721 b, type->storage_class, type->deref, NULL);
1722 nir_address_format addr_format = vtn_mode_to_address_format(b, mode);
1723
1724 const nir_const_value *null_value = nir_address_format_null_value(addr_format);
1725 memcpy(c->values, null_value,
1726 sizeof(nir_const_value) * nir_address_format_num_components(addr_format));
1727 break;
1728 }
1729
1730 case vtn_base_type_void:
1731 case vtn_base_type_image:
1732 case vtn_base_type_sampler:
1733 case vtn_base_type_sampled_image:
1734 case vtn_base_type_function:
1735 /* For those we have to return something but it doesn't matter what. */
1736 break;
1737
1738 case vtn_base_type_matrix:
1739 case vtn_base_type_array:
1740 vtn_assert(type->length > 0);
1741 c->num_elements = type->length;
1742 c->elements = ralloc_array(b, nir_constant *, c->num_elements);
1743
1744 c->elements[0] = vtn_null_constant(b, type->array_element);
1745 for (unsigned i = 1; i < c->num_elements; i++)
1746 c->elements[i] = c->elements[0];
1747 break;
1748
1749 case vtn_base_type_struct:
1750 c->num_elements = type->length;
1751 c->elements = ralloc_array(b, nir_constant *, c->num_elements);
1752 for (unsigned i = 0; i < c->num_elements; i++)
1753 c->elements[i] = vtn_null_constant(b, type->members[i]);
1754 break;
1755
1756 default:
1757 vtn_fail("Invalid type for null constant");
1758 }
1759
1760 return c;
1761 }
1762
1763 static void
1764 spec_constant_decoration_cb(struct vtn_builder *b, UNUSED struct vtn_value *val,
1765 ASSERTED int member,
1766 const struct vtn_decoration *dec, void *data)
1767 {
1768 vtn_assert(member == -1);
1769 if (dec->decoration != SpvDecorationSpecId)
1770 return;
1771
1772 nir_const_value *value = data;
1773 for (unsigned i = 0; i < b->num_specializations; i++) {
1774 if (b->specializations[i].id == dec->operands[0]) {
1775 *value = b->specializations[i].value;
1776 return;
1777 }
1778 }
1779 }
1780
1781 static void
1782 handle_workgroup_size_decoration_cb(struct vtn_builder *b,
1783 struct vtn_value *val,
1784 ASSERTED int member,
1785 const struct vtn_decoration *dec,
1786 UNUSED void *data)
1787 {
1788 vtn_assert(member == -1);
1789 if (dec->decoration != SpvDecorationBuiltIn ||
1790 dec->operands[0] != SpvBuiltInWorkgroupSize)
1791 return;
1792
1793 vtn_assert(val->type->type == glsl_vector_type(GLSL_TYPE_UINT, 3));
1794 b->workgroup_size_builtin = val;
1795 }
1796
1797 static void
1798 vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
1799 const uint32_t *w, unsigned count)
1800 {
1801 struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_constant);
1802 val->constant = rzalloc(b, nir_constant);
1803 switch (opcode) {
1804 case SpvOpConstantTrue:
1805 case SpvOpConstantFalse:
1806 case SpvOpSpecConstantTrue:
1807 case SpvOpSpecConstantFalse: {
1808 vtn_fail_if(val->type->type != glsl_bool_type(),
1809 "Result type of %s must be OpTypeBool",
1810 spirv_op_to_string(opcode));
1811
1812 bool bval = (opcode == SpvOpConstantTrue ||
1813 opcode == SpvOpSpecConstantTrue);
1814
1815 nir_const_value u32val = nir_const_value_for_uint(bval, 32);
1816
1817 if (opcode == SpvOpSpecConstantTrue ||
1818 opcode == SpvOpSpecConstantFalse)
1819 vtn_foreach_decoration(b, val, spec_constant_decoration_cb, &u32val);
1820
1821 val->constant->values[0].b = u32val.u32 != 0;
1822 break;
1823 }
1824
1825 case SpvOpConstant:
1826 case SpvOpSpecConstant: {
1827 vtn_fail_if(val->type->base_type != vtn_base_type_scalar,
1828 "Result type of %s must be a scalar",
1829 spirv_op_to_string(opcode));
1830 int bit_size = glsl_get_bit_size(val->type->type);
1831 switch (bit_size) {
1832 case 64:
1833 val->constant->values[0].u64 = vtn_u64_literal(&w[3]);
1834 break;
1835 case 32:
1836 val->constant->values[0].u32 = w[3];
1837 break;
1838 case 16:
1839 val->constant->values[0].u16 = w[3];
1840 break;
1841 case 8:
1842 val->constant->values[0].u8 = w[3];
1843 break;
1844 default:
1845 vtn_fail("Unsupported SpvOpConstant bit size: %u", bit_size);
1846 }
1847
1848 if (opcode == SpvOpSpecConstant)
1849 vtn_foreach_decoration(b, val, spec_constant_decoration_cb,
1850 &val->constant->values[0]);
1851 break;
1852 }
1853
1854 case SpvOpSpecConstantComposite:
1855 case SpvOpConstantComposite: {
1856 unsigned elem_count = count - 3;
1857 vtn_fail_if(elem_count != val->type->length,
1858 "%s has %u constituents, expected %u",
1859 spirv_op_to_string(opcode), elem_count, val->type->length);
1860
1861 nir_constant **elems = ralloc_array(b, nir_constant *, elem_count);
1862 for (unsigned i = 0; i < elem_count; i++) {
1863 struct vtn_value *val = vtn_untyped_value(b, w[i + 3]);
1864
1865 if (val->value_type == vtn_value_type_constant) {
1866 elems[i] = val->constant;
1867 } else {
1868 vtn_fail_if(val->value_type != vtn_value_type_undef,
1869 "only constants or undefs allowed for "
1870 "SpvOpConstantComposite");
1871 /* to make it easier, just insert a NULL constant for now */
1872 elems[i] = vtn_null_constant(b, val->type);
1873 }
1874 }
1875
1876 switch (val->type->base_type) {
1877 case vtn_base_type_vector: {
1878 assert(glsl_type_is_vector(val->type->type));
1879 for (unsigned i = 0; i < elem_count; i++)
1880 val->constant->values[i] = elems[i]->values[0];
1881 break;
1882 }
1883
1884 case vtn_base_type_matrix:
1885 case vtn_base_type_struct:
1886 case vtn_base_type_array:
1887 ralloc_steal(val->constant, elems);
1888 val->constant->num_elements = elem_count;
1889 val->constant->elements = elems;
1890 break;
1891
1892 default:
1893 vtn_fail("Result type of %s must be a composite type",
1894 spirv_op_to_string(opcode));
1895 }
1896 break;
1897 }
1898
1899 case SpvOpSpecConstantOp: {
1900 nir_const_value u32op = nir_const_value_for_uint(w[3], 32);
1901 vtn_foreach_decoration(b, val, spec_constant_decoration_cb, &u32op);
1902 SpvOp opcode = u32op.u32;
1903 switch (opcode) {
1904 case SpvOpVectorShuffle: {
1905 struct vtn_value *v0 = &b->values[w[4]];
1906 struct vtn_value *v1 = &b->values[w[5]];
1907
1908 vtn_assert(v0->value_type == vtn_value_type_constant ||
1909 v0->value_type == vtn_value_type_undef);
1910 vtn_assert(v1->value_type == vtn_value_type_constant ||
1911 v1->value_type == vtn_value_type_undef);
1912
1913 unsigned len0 = glsl_get_vector_elements(v0->type->type);
1914 unsigned len1 = glsl_get_vector_elements(v1->type->type);
1915
1916 vtn_assert(len0 + len1 < 16);
1917
1918 unsigned bit_size = glsl_get_bit_size(val->type->type);
1919 unsigned bit_size0 = glsl_get_bit_size(v0->type->type);
1920 unsigned bit_size1 = glsl_get_bit_size(v1->type->type);
1921
1922 vtn_assert(bit_size == bit_size0 && bit_size == bit_size1);
1923 (void)bit_size0; (void)bit_size1;
1924
1925 nir_const_value undef = { .u64 = 0xdeadbeefdeadbeef };
1926 nir_const_value combined[NIR_MAX_VEC_COMPONENTS * 2];
1927
1928 if (v0->value_type == vtn_value_type_constant) {
1929 for (unsigned i = 0; i < len0; i++)
1930 combined[i] = v0->constant->values[i];
1931 }
1932 if (v1->value_type == vtn_value_type_constant) {
1933 for (unsigned i = 0; i < len1; i++)
1934 combined[len0 + i] = v1->constant->values[i];
1935 }
1936
1937 for (unsigned i = 0, j = 0; i < count - 6; i++, j++) {
1938 uint32_t comp = w[i + 6];
1939 if (comp == (uint32_t)-1) {
1940 /* If component is not used, set the value to a known constant
1941 * to detect if it is wrongly used.
1942 */
1943 val->constant->values[j] = undef;
1944 } else {
1945 vtn_fail_if(comp >= len0 + len1,
1946 "All Component literals must either be FFFFFFFF "
1947 "or in [0, N - 1] (inclusive).");
1948 val->constant->values[j] = combined[comp];
1949 }
1950 }
1951 break;
1952 }
1953
1954 case SpvOpCompositeExtract:
1955 case SpvOpCompositeInsert: {
1956 struct vtn_value *comp;
1957 unsigned deref_start;
1958 struct nir_constant **c;
1959 if (opcode == SpvOpCompositeExtract) {
1960 comp = vtn_value(b, w[4], vtn_value_type_constant);
1961 deref_start = 5;
1962 c = &comp->constant;
1963 } else {
1964 comp = vtn_value(b, w[5], vtn_value_type_constant);
1965 deref_start = 6;
1966 val->constant = nir_constant_clone(comp->constant,
1967 (nir_variable *)b);
1968 c = &val->constant;
1969 }
1970
1971 int elem = -1;
1972 const struct vtn_type *type = comp->type;
1973 for (unsigned i = deref_start; i < count; i++) {
1974 vtn_fail_if(w[i] > type->length,
1975 "%uth index of %s is %u but the type has only "
1976 "%u elements", i - deref_start,
1977 spirv_op_to_string(opcode), w[i], type->length);
1978
1979 switch (type->base_type) {
1980 case vtn_base_type_vector:
1981 elem = w[i];
1982 type = type->array_element;
1983 break;
1984
1985 case vtn_base_type_matrix:
1986 case vtn_base_type_array:
1987 c = &(*c)->elements[w[i]];
1988 type = type->array_element;
1989 break;
1990
1991 case vtn_base_type_struct:
1992 c = &(*c)->elements[w[i]];
1993 type = type->members[w[i]];
1994 break;
1995
1996 default:
1997 vtn_fail("%s must only index into composite types",
1998 spirv_op_to_string(opcode));
1999 }
2000 }
2001
2002 if (opcode == SpvOpCompositeExtract) {
2003 if (elem == -1) {
2004 val->constant = *c;
2005 } else {
2006 unsigned num_components = type->length;
2007 for (unsigned i = 0; i < num_components; i++)
2008 val->constant->values[i] = (*c)->values[elem + i];
2009 }
2010 } else {
2011 struct vtn_value *insert =
2012 vtn_value(b, w[4], vtn_value_type_constant);
2013 vtn_assert(insert->type == type);
2014 if (elem == -1) {
2015 *c = insert->constant;
2016 } else {
2017 unsigned num_components = type->length;
2018 for (unsigned i = 0; i < num_components; i++)
2019 (*c)->values[elem + i] = insert->constant->values[i];
2020 }
2021 }
2022 break;
2023 }
2024
2025 default: {
2026 bool swap;
2027 nir_alu_type dst_alu_type = nir_get_nir_type_for_glsl_type(val->type->type);
2028 nir_alu_type src_alu_type = dst_alu_type;
2029 unsigned num_components = glsl_get_vector_elements(val->type->type);
2030 unsigned bit_size;
2031
2032 vtn_assert(count <= 7);
2033
2034 switch (opcode) {
2035 case SpvOpSConvert:
2036 case SpvOpFConvert:
2037 case SpvOpUConvert:
2038 /* We have a source in a conversion */
2039 src_alu_type =
2040 nir_get_nir_type_for_glsl_type(vtn_get_value_type(b, w[4])->type);
2041 /* We use the bitsize of the conversion source to evaluate the opcode later */
2042 bit_size = glsl_get_bit_size(vtn_get_value_type(b, w[4])->type);
2043 break;
2044 default:
2045 bit_size = glsl_get_bit_size(val->type->type);
2046 };
2047
2048 nir_op op = vtn_nir_alu_op_for_spirv_opcode(b, opcode, &swap,
2049 nir_alu_type_get_type_size(src_alu_type),
2050 nir_alu_type_get_type_size(dst_alu_type));
2051 nir_const_value src[3][NIR_MAX_VEC_COMPONENTS];
2052
2053 for (unsigned i = 0; i < count - 4; i++) {
2054 struct vtn_value *src_val =
2055 vtn_value(b, w[4 + i], vtn_value_type_constant);
2056
2057 /* If this is an unsized source, pull the bit size from the
2058 * source; otherwise, we'll use the bit size from the destination.
2059 */
2060 if (!nir_alu_type_get_type_size(nir_op_infos[op].input_types[i]))
2061 bit_size = glsl_get_bit_size(src_val->type->type);
2062
2063 unsigned src_comps = nir_op_infos[op].input_sizes[i] ?
2064 nir_op_infos[op].input_sizes[i] :
2065 num_components;
2066
2067 unsigned j = swap ? 1 - i : i;
2068 for (unsigned c = 0; c < src_comps; c++)
2069 src[j][c] = src_val->constant->values[c];
2070 }
2071
2072 /* fix up fixed size sources */
2073 switch (op) {
2074 case nir_op_ishl:
2075 case nir_op_ishr:
2076 case nir_op_ushr: {
2077 if (bit_size == 32)
2078 break;
2079 for (unsigned i = 0; i < num_components; ++i) {
2080 switch (bit_size) {
2081 case 64: src[1][i].u32 = src[1][i].u64; break;
2082 case 16: src[1][i].u32 = src[1][i].u16; break;
2083 case 8: src[1][i].u32 = src[1][i].u8; break;
2084 }
2085 }
2086 break;
2087 }
2088 default:
2089 break;
2090 }
2091
2092 nir_const_value *srcs[3] = {
2093 src[0], src[1], src[2],
2094 };
2095 nir_eval_const_opcode(op, val->constant->values,
2096 num_components, bit_size, srcs,
2097 b->shader->info.float_controls_execution_mode);
2098 break;
2099 } /* default */
2100 }
2101 break;
2102 }
2103
2104 case SpvOpConstantNull:
2105 val->constant = vtn_null_constant(b, val->type);
2106 break;
2107
2108 default:
2109 vtn_fail_with_opcode("Unhandled opcode", opcode);
2110 }
2111
2112 /* Now that we have the value, update the workgroup size if needed */
2113 vtn_foreach_decoration(b, val, handle_workgroup_size_decoration_cb, NULL);
2114 }
2115
2116 static void
2117 vtn_split_barrier_semantics(struct vtn_builder *b,
2118 SpvMemorySemanticsMask semantics,
2119 SpvMemorySemanticsMask *before,
2120 SpvMemorySemanticsMask *after)
2121 {
2122 /* For memory semantics embedded in operations, we split them into up to
2123 * two barriers, to be added before and after the operation. This is less
2124 * strict than if we propagated until the final backend stage, but still
2125 * result in correct execution.
2126 *
2127 * A further improvement could be pipe this information (and use!) into the
2128 * next compiler layers, at the expense of making the handling of barriers
2129 * more complicated.
2130 */
2131
2132 *before = SpvMemorySemanticsMaskNone;
2133 *after = SpvMemorySemanticsMaskNone;
2134
2135 SpvMemorySemanticsMask order_semantics =
2136 semantics & (SpvMemorySemanticsAcquireMask |
2137 SpvMemorySemanticsReleaseMask |
2138 SpvMemorySemanticsAcquireReleaseMask |
2139 SpvMemorySemanticsSequentiallyConsistentMask);
2140
2141 if (util_bitcount(order_semantics) > 1) {
2142 /* Old GLSLang versions incorrectly set all the ordering bits. This was
2143 * fixed in c51287d744fb6e7e9ccc09f6f8451e6c64b1dad6 of glslang repo,
2144 * and it is in GLSLang since revision "SPIRV99.1321" (from Jul-2016).
2145 */
2146 vtn_warn("Multiple memory ordering semantics specified, "
2147 "assuming AcquireRelease.");
2148 order_semantics = SpvMemorySemanticsAcquireReleaseMask;
2149 }
2150
2151 const SpvMemorySemanticsMask av_vis_semantics =
2152 semantics & (SpvMemorySemanticsMakeAvailableMask |
2153 SpvMemorySemanticsMakeVisibleMask);
2154
2155 const SpvMemorySemanticsMask storage_semantics =
2156 semantics & (SpvMemorySemanticsUniformMemoryMask |
2157 SpvMemorySemanticsSubgroupMemoryMask |
2158 SpvMemorySemanticsWorkgroupMemoryMask |
2159 SpvMemorySemanticsCrossWorkgroupMemoryMask |
2160 SpvMemorySemanticsAtomicCounterMemoryMask |
2161 SpvMemorySemanticsImageMemoryMask |
2162 SpvMemorySemanticsOutputMemoryMask);
2163
2164 const SpvMemorySemanticsMask other_semantics =
2165 semantics & ~(order_semantics | av_vis_semantics | storage_semantics |
2166 SpvMemorySemanticsVolatileMask);
2167
2168 if (other_semantics)
2169 vtn_warn("Ignoring unhandled memory semantics: %u\n", other_semantics);
2170
2171 /* SequentiallyConsistent is treated as AcquireRelease. */
2172
2173 /* The RELEASE barrier happens BEFORE the operation, and it is usually
2174 * associated with a Store. All the write operations with a matching
2175 * semantics will not be reordered after the Store.
2176 */
2177 if (order_semantics & (SpvMemorySemanticsReleaseMask |
2178 SpvMemorySemanticsAcquireReleaseMask |
2179 SpvMemorySemanticsSequentiallyConsistentMask)) {
2180 *before |= SpvMemorySemanticsReleaseMask | storage_semantics;
2181 }
2182
2183 /* The ACQUIRE barrier happens AFTER the operation, and it is usually
2184 * associated with a Load. All the operations with a matching semantics
2185 * will not be reordered before the Load.
2186 */
2187 if (order_semantics & (SpvMemorySemanticsAcquireMask |
2188 SpvMemorySemanticsAcquireReleaseMask |
2189 SpvMemorySemanticsSequentiallyConsistentMask)) {
2190 *after |= SpvMemorySemanticsAcquireMask | storage_semantics;
2191 }
2192
2193 if (av_vis_semantics & SpvMemorySemanticsMakeVisibleMask)
2194 *before |= SpvMemorySemanticsMakeVisibleMask | storage_semantics;
2195
2196 if (av_vis_semantics & SpvMemorySemanticsMakeAvailableMask)
2197 *after |= SpvMemorySemanticsMakeAvailableMask | storage_semantics;
2198 }
2199
2200 static nir_memory_semantics
2201 vtn_mem_semantics_to_nir_mem_semantics(struct vtn_builder *b,
2202 SpvMemorySemanticsMask semantics)
2203 {
2204 nir_memory_semantics nir_semantics = 0;
2205
2206 SpvMemorySemanticsMask order_semantics =
2207 semantics & (SpvMemorySemanticsAcquireMask |
2208 SpvMemorySemanticsReleaseMask |
2209 SpvMemorySemanticsAcquireReleaseMask |
2210 SpvMemorySemanticsSequentiallyConsistentMask);
2211
2212 if (util_bitcount(order_semantics) > 1) {
2213 /* Old GLSLang versions incorrectly set all the ordering bits. This was
2214 * fixed in c51287d744fb6e7e9ccc09f6f8451e6c64b1dad6 of glslang repo,
2215 * and it is in GLSLang since revision "SPIRV99.1321" (from Jul-2016).
2216 */
2217 vtn_warn("Multiple memory ordering semantics bits specified, "
2218 "assuming AcquireRelease.");
2219 order_semantics = SpvMemorySemanticsAcquireReleaseMask;
2220 }
2221
2222 switch (order_semantics) {
2223 case 0:
2224 /* Not an ordering barrier. */
2225 break;
2226
2227 case SpvMemorySemanticsAcquireMask:
2228 nir_semantics = NIR_MEMORY_ACQUIRE;
2229 break;
2230
2231 case SpvMemorySemanticsReleaseMask:
2232 nir_semantics = NIR_MEMORY_RELEASE;
2233 break;
2234
2235 case SpvMemorySemanticsSequentiallyConsistentMask:
2236 /* Fall through. Treated as AcquireRelease in Vulkan. */
2237 case SpvMemorySemanticsAcquireReleaseMask:
2238 nir_semantics = NIR_MEMORY_ACQUIRE | NIR_MEMORY_RELEASE;
2239 break;
2240
2241 default:
2242 unreachable("Invalid memory order semantics");
2243 }
2244
2245 if (semantics & SpvMemorySemanticsMakeAvailableMask) {
2246 vtn_fail_if(!b->options->caps.vk_memory_model,
2247 "To use MakeAvailable memory semantics the VulkanMemoryModel "
2248 "capability must be declared.");
2249 nir_semantics |= NIR_MEMORY_MAKE_AVAILABLE;
2250 }
2251
2252 if (semantics & SpvMemorySemanticsMakeVisibleMask) {
2253 vtn_fail_if(!b->options->caps.vk_memory_model,
2254 "To use MakeVisible memory semantics the VulkanMemoryModel "
2255 "capability must be declared.");
2256 nir_semantics |= NIR_MEMORY_MAKE_VISIBLE;
2257 }
2258
2259 return nir_semantics;
2260 }
2261
2262 static nir_variable_mode
2263 vtn_mem_sematics_to_nir_var_modes(struct vtn_builder *b,
2264 SpvMemorySemanticsMask semantics)
2265 {
2266 /* Vulkan Environment for SPIR-V says "SubgroupMemory, CrossWorkgroupMemory,
2267 * and AtomicCounterMemory are ignored".
2268 */
2269 semantics &= ~(SpvMemorySemanticsSubgroupMemoryMask |
2270 SpvMemorySemanticsCrossWorkgroupMemoryMask |
2271 SpvMemorySemanticsAtomicCounterMemoryMask);
2272
2273 /* TODO: Consider adding nir_var_mem_image mode to NIR so it can be used
2274 * for SpvMemorySemanticsImageMemoryMask.
2275 */
2276
2277 nir_variable_mode modes = 0;
2278 if (semantics & (SpvMemorySemanticsUniformMemoryMask |
2279 SpvMemorySemanticsImageMemoryMask)) {
2280 modes |= nir_var_uniform |
2281 nir_var_mem_ubo |
2282 nir_var_mem_ssbo |
2283 nir_var_mem_global;
2284 }
2285 if (semantics & SpvMemorySemanticsWorkgroupMemoryMask)
2286 modes |= nir_var_mem_shared;
2287 if (semantics & SpvMemorySemanticsOutputMemoryMask) {
2288 modes |= nir_var_shader_out;
2289 }
2290
2291 return modes;
2292 }
2293
2294 static nir_scope
2295 vtn_scope_to_nir_scope(struct vtn_builder *b, SpvScope scope)
2296 {
2297 nir_scope nir_scope;
2298 switch (scope) {
2299 case SpvScopeDevice:
2300 vtn_fail_if(b->options->caps.vk_memory_model &&
2301 !b->options->caps.vk_memory_model_device_scope,
2302 "If the Vulkan memory model is declared and any instruction "
2303 "uses Device scope, the VulkanMemoryModelDeviceScope "
2304 "capability must be declared.");
2305 nir_scope = NIR_SCOPE_DEVICE;
2306 break;
2307
2308 case SpvScopeQueueFamily:
2309 vtn_fail_if(!b->options->caps.vk_memory_model,
2310 "To use Queue Family scope, the VulkanMemoryModel capability "
2311 "must be declared.");
2312 nir_scope = NIR_SCOPE_QUEUE_FAMILY;
2313 break;
2314
2315 case SpvScopeWorkgroup:
2316 nir_scope = NIR_SCOPE_WORKGROUP;
2317 break;
2318
2319 case SpvScopeSubgroup:
2320 nir_scope = NIR_SCOPE_SUBGROUP;
2321 break;
2322
2323 case SpvScopeInvocation:
2324 nir_scope = NIR_SCOPE_INVOCATION;
2325 break;
2326
2327 default:
2328 vtn_fail("Invalid memory scope");
2329 }
2330
2331 return nir_scope;
2332 }
2333
2334 static void
2335 vtn_emit_scoped_control_barrier(struct vtn_builder *b, SpvScope exec_scope,
2336 SpvScope mem_scope,
2337 SpvMemorySemanticsMask semantics)
2338 {
2339 nir_memory_semantics nir_semantics =
2340 vtn_mem_semantics_to_nir_mem_semantics(b, semantics);
2341 nir_variable_mode modes = vtn_mem_sematics_to_nir_var_modes(b, semantics);
2342 nir_scope nir_exec_scope = vtn_scope_to_nir_scope(b, exec_scope);
2343
2344 /* Memory semantics is optional for OpControlBarrier. */
2345 nir_scope nir_mem_scope;
2346 if (nir_semantics == 0 || modes == 0)
2347 nir_mem_scope = NIR_SCOPE_NONE;
2348 else
2349 nir_mem_scope = vtn_scope_to_nir_scope(b, mem_scope);
2350
2351 nir_scoped_barrier(&b->nb, nir_exec_scope, nir_mem_scope, nir_semantics, modes);
2352 }
2353
2354 static void
2355 vtn_emit_scoped_memory_barrier(struct vtn_builder *b, SpvScope scope,
2356 SpvMemorySemanticsMask semantics)
2357 {
2358 nir_variable_mode modes = vtn_mem_sematics_to_nir_var_modes(b, semantics);
2359 nir_memory_semantics nir_semantics =
2360 vtn_mem_semantics_to_nir_mem_semantics(b, semantics);
2361
2362 /* No barrier to add. */
2363 if (nir_semantics == 0 || modes == 0)
2364 return;
2365
2366 nir_scope nir_mem_scope = vtn_scope_to_nir_scope(b, scope);
2367 nir_scoped_barrier(&b->nb, NIR_SCOPE_NONE, nir_mem_scope, nir_semantics, modes);
2368 }
2369
2370 struct vtn_ssa_value *
2371 vtn_create_ssa_value(struct vtn_builder *b, const struct glsl_type *type)
2372 {
2373 /* Always use bare types for SSA values for a couple of reasons:
2374 *
2375 * 1. Code which emits deref chains should never listen to the explicit
2376 * layout information on the SSA value if any exists. If we've
2377 * accidentally been relying on this, we want to find those bugs.
2378 *
2379 * 2. We want to be able to quickly check that an SSA value being assigned
2380 * to a SPIR-V value has the right type. Using bare types everywhere
2381 * ensures that we can pointer-compare.
2382 */
2383 struct vtn_ssa_value *val = rzalloc(b, struct vtn_ssa_value);
2384 val->type = glsl_get_bare_type(type);
2385
2386
2387 if (!glsl_type_is_vector_or_scalar(type)) {
2388 unsigned elems = glsl_get_length(val->type);
2389 val->elems = ralloc_array(b, struct vtn_ssa_value *, elems);
2390 if (glsl_type_is_array_or_matrix(type)) {
2391 const struct glsl_type *elem_type = glsl_get_array_element(type);
2392 for (unsigned i = 0; i < elems; i++)
2393 val->elems[i] = vtn_create_ssa_value(b, elem_type);
2394 } else {
2395 vtn_assert(glsl_type_is_struct_or_ifc(type));
2396 for (unsigned i = 0; i < elems; i++) {
2397 const struct glsl_type *elem_type = glsl_get_struct_field(type, i);
2398 val->elems[i] = vtn_create_ssa_value(b, elem_type);
2399 }
2400 }
2401 }
2402
2403 return val;
2404 }
2405
2406 static nir_tex_src
2407 vtn_tex_src(struct vtn_builder *b, unsigned index, nir_tex_src_type type)
2408 {
2409 nir_tex_src src;
2410 src.src = nir_src_for_ssa(vtn_get_nir_ssa(b, index));
2411 src.src_type = type;
2412 return src;
2413 }
2414
2415 static uint32_t
2416 image_operand_arg(struct vtn_builder *b, const uint32_t *w, uint32_t count,
2417 uint32_t mask_idx, SpvImageOperandsMask op)
2418 {
2419 static const SpvImageOperandsMask ops_with_arg =
2420 SpvImageOperandsBiasMask |
2421 SpvImageOperandsLodMask |
2422 SpvImageOperandsGradMask |
2423 SpvImageOperandsConstOffsetMask |
2424 SpvImageOperandsOffsetMask |
2425 SpvImageOperandsConstOffsetsMask |
2426 SpvImageOperandsSampleMask |
2427 SpvImageOperandsMinLodMask |
2428 SpvImageOperandsMakeTexelAvailableMask |
2429 SpvImageOperandsMakeTexelVisibleMask;
2430
2431 assert(util_bitcount(op) == 1);
2432 assert(w[mask_idx] & op);
2433 assert(op & ops_with_arg);
2434
2435 uint32_t idx = util_bitcount(w[mask_idx] & (op - 1) & ops_with_arg) + 1;
2436
2437 /* Adjust indices for operands with two arguments. */
2438 static const SpvImageOperandsMask ops_with_two_args =
2439 SpvImageOperandsGradMask;
2440 idx += util_bitcount(w[mask_idx] & (op - 1) & ops_with_two_args);
2441
2442 idx += mask_idx;
2443
2444 vtn_fail_if(idx + (op & ops_with_two_args ? 1 : 0) >= count,
2445 "Image op claims to have %s but does not enough "
2446 "following operands", spirv_imageoperands_to_string(op));
2447
2448 return idx;
2449 }
2450
2451 static void
2452 non_uniform_decoration_cb(struct vtn_builder *b,
2453 struct vtn_value *val, int member,
2454 const struct vtn_decoration *dec, void *void_ctx)
2455 {
2456 enum gl_access_qualifier *access = void_ctx;
2457 switch (dec->decoration) {
2458 case SpvDecorationNonUniformEXT:
2459 *access |= ACCESS_NON_UNIFORM;
2460 break;
2461
2462 default:
2463 break;
2464 }
2465 }
2466
2467 static void
2468 vtn_handle_texture(struct vtn_builder *b, SpvOp opcode,
2469 const uint32_t *w, unsigned count)
2470 {
2471 struct vtn_type *ret_type = vtn_get_type(b, w[1]);
2472
2473 if (opcode == SpvOpSampledImage) {
2474 struct vtn_sampled_image si = {
2475 .image = vtn_get_image(b, w[3]),
2476 .sampler = vtn_get_sampler(b, w[4]),
2477 };
2478
2479 enum gl_access_qualifier access = 0;
2480 vtn_foreach_decoration(b, vtn_untyped_value(b, w[3]),
2481 non_uniform_decoration_cb, &access);
2482 vtn_foreach_decoration(b, vtn_untyped_value(b, w[4]),
2483 non_uniform_decoration_cb, &access);
2484
2485 vtn_push_sampled_image(b, w[2], si, access & ACCESS_NON_UNIFORM);
2486 return;
2487 } else if (opcode == SpvOpImage) {
2488 struct vtn_sampled_image si = vtn_get_sampled_image(b, w[3]);
2489
2490 enum gl_access_qualifier access = 0;
2491 vtn_foreach_decoration(b, vtn_untyped_value(b, w[3]),
2492 non_uniform_decoration_cb, &access);
2493
2494 vtn_push_image(b, w[2], si.image, access & ACCESS_NON_UNIFORM);
2495 return;
2496 }
2497
2498 nir_deref_instr *image = NULL, *sampler = NULL;
2499 struct vtn_value *sampled_val = vtn_untyped_value(b, w[3]);
2500 if (sampled_val->type->base_type == vtn_base_type_sampled_image) {
2501 struct vtn_sampled_image si = vtn_get_sampled_image(b, w[3]);
2502 image = si.image;
2503 sampler = si.sampler;
2504 } else {
2505 image = vtn_get_image(b, w[3]);
2506 }
2507
2508 const enum glsl_sampler_dim sampler_dim = glsl_get_sampler_dim(image->type);
2509 const bool is_array = glsl_sampler_type_is_array(image->type);
2510 nir_alu_type dest_type = nir_type_invalid;
2511
2512 /* Figure out the base texture operation */
2513 nir_texop texop;
2514 switch (opcode) {
2515 case SpvOpImageSampleImplicitLod:
2516 case SpvOpImageSampleDrefImplicitLod:
2517 case SpvOpImageSampleProjImplicitLod:
2518 case SpvOpImageSampleProjDrefImplicitLod:
2519 texop = nir_texop_tex;
2520 break;
2521
2522 case SpvOpImageSampleExplicitLod:
2523 case SpvOpImageSampleDrefExplicitLod:
2524 case SpvOpImageSampleProjExplicitLod:
2525 case SpvOpImageSampleProjDrefExplicitLod:
2526 texop = nir_texop_txl;
2527 break;
2528
2529 case SpvOpImageFetch:
2530 if (sampler_dim == GLSL_SAMPLER_DIM_MS) {
2531 texop = nir_texop_txf_ms;
2532 } else {
2533 texop = nir_texop_txf;
2534 }
2535 break;
2536
2537 case SpvOpImageGather:
2538 case SpvOpImageDrefGather:
2539 texop = nir_texop_tg4;
2540 break;
2541
2542 case SpvOpImageQuerySizeLod:
2543 case SpvOpImageQuerySize:
2544 texop = nir_texop_txs;
2545 dest_type = nir_type_int;
2546 break;
2547
2548 case SpvOpImageQueryLod:
2549 texop = nir_texop_lod;
2550 dest_type = nir_type_float;
2551 break;
2552
2553 case SpvOpImageQueryLevels:
2554 texop = nir_texop_query_levels;
2555 dest_type = nir_type_int;
2556 break;
2557
2558 case SpvOpImageQuerySamples:
2559 texop = nir_texop_texture_samples;
2560 dest_type = nir_type_int;
2561 break;
2562
2563 case SpvOpFragmentFetchAMD:
2564 texop = nir_texop_fragment_fetch;
2565 break;
2566
2567 case SpvOpFragmentMaskFetchAMD:
2568 texop = nir_texop_fragment_mask_fetch;
2569 dest_type = nir_type_uint;
2570 break;
2571
2572 default:
2573 vtn_fail_with_opcode("Unhandled opcode", opcode);
2574 }
2575
2576 nir_tex_src srcs[10]; /* 10 should be enough */
2577 nir_tex_src *p = srcs;
2578
2579 p->src = nir_src_for_ssa(&image->dest.ssa);
2580 p->src_type = nir_tex_src_texture_deref;
2581 p++;
2582
2583 switch (texop) {
2584 case nir_texop_tex:
2585 case nir_texop_txb:
2586 case nir_texop_txl:
2587 case nir_texop_txd:
2588 case nir_texop_tg4:
2589 case nir_texop_lod:
2590 vtn_fail_if(sampler == NULL,
2591 "%s requires an image of type OpTypeSampledImage",
2592 spirv_op_to_string(opcode));
2593 p->src = nir_src_for_ssa(&sampler->dest.ssa);
2594 p->src_type = nir_tex_src_sampler_deref;
2595 p++;
2596 break;
2597 case nir_texop_txf:
2598 case nir_texop_txf_ms:
2599 case nir_texop_txs:
2600 case nir_texop_query_levels:
2601 case nir_texop_texture_samples:
2602 case nir_texop_samples_identical:
2603 case nir_texop_fragment_fetch:
2604 case nir_texop_fragment_mask_fetch:
2605 /* These don't */
2606 break;
2607 case nir_texop_txf_ms_fb:
2608 vtn_fail("unexpected nir_texop_txf_ms_fb");
2609 break;
2610 case nir_texop_txf_ms_mcs:
2611 vtn_fail("unexpected nir_texop_txf_ms_mcs");
2612 case nir_texop_tex_prefetch:
2613 vtn_fail("unexpected nir_texop_tex_prefetch");
2614 }
2615
2616 unsigned idx = 4;
2617
2618 struct nir_ssa_def *coord;
2619 unsigned coord_components;
2620 switch (opcode) {
2621 case SpvOpImageSampleImplicitLod:
2622 case SpvOpImageSampleExplicitLod:
2623 case SpvOpImageSampleDrefImplicitLod:
2624 case SpvOpImageSampleDrefExplicitLod:
2625 case SpvOpImageSampleProjImplicitLod:
2626 case SpvOpImageSampleProjExplicitLod:
2627 case SpvOpImageSampleProjDrefImplicitLod:
2628 case SpvOpImageSampleProjDrefExplicitLod:
2629 case SpvOpImageFetch:
2630 case SpvOpImageGather:
2631 case SpvOpImageDrefGather:
2632 case SpvOpImageQueryLod:
2633 case SpvOpFragmentFetchAMD:
2634 case SpvOpFragmentMaskFetchAMD: {
2635 /* All these types have the coordinate as their first real argument */
2636 coord_components = glsl_get_sampler_dim_coordinate_components(sampler_dim);
2637
2638 if (is_array && texop != nir_texop_lod)
2639 coord_components++;
2640
2641 struct vtn_ssa_value *coord_val = vtn_ssa_value(b, w[idx++]);
2642 coord = coord_val->def;
2643 p->src = nir_src_for_ssa(nir_channels(&b->nb, coord,
2644 (1 << coord_components) - 1));
2645
2646 /* OpenCL allows integer sampling coordinates */
2647 if (glsl_type_is_integer(coord_val->type) &&
2648 opcode == SpvOpImageSampleExplicitLod) {
2649 vtn_fail_if(b->shader->info.stage != MESA_SHADER_KERNEL,
2650 "Unless the Kernel capability is being used, the coordinate parameter "
2651 "OpImageSampleExplicitLod must be floating point.");
2652
2653 p->src = nir_src_for_ssa(nir_i2f32(&b->nb, p->src.ssa));
2654 }
2655
2656 p->src_type = nir_tex_src_coord;
2657 p++;
2658 break;
2659 }
2660
2661 default:
2662 coord = NULL;
2663 coord_components = 0;
2664 break;
2665 }
2666
2667 switch (opcode) {
2668 case SpvOpImageSampleProjImplicitLod:
2669 case SpvOpImageSampleProjExplicitLod:
2670 case SpvOpImageSampleProjDrefImplicitLod:
2671 case SpvOpImageSampleProjDrefExplicitLod:
2672 /* These have the projector as the last coordinate component */
2673 p->src = nir_src_for_ssa(nir_channel(&b->nb, coord, coord_components));
2674 p->src_type = nir_tex_src_projector;
2675 p++;
2676 break;
2677
2678 default:
2679 break;
2680 }
2681
2682 bool is_shadow = false;
2683 unsigned gather_component = 0;
2684 switch (opcode) {
2685 case SpvOpImageSampleDrefImplicitLod:
2686 case SpvOpImageSampleDrefExplicitLod:
2687 case SpvOpImageSampleProjDrefImplicitLod:
2688 case SpvOpImageSampleProjDrefExplicitLod:
2689 case SpvOpImageDrefGather:
2690 /* These all have an explicit depth value as their next source */
2691 is_shadow = true;
2692 (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_comparator);
2693 break;
2694
2695 case SpvOpImageGather:
2696 /* This has a component as its next source */
2697 gather_component = vtn_constant_uint(b, w[idx++]);
2698 break;
2699
2700 default:
2701 break;
2702 }
2703
2704 /* For OpImageQuerySizeLod, we always have an LOD */
2705 if (opcode == SpvOpImageQuerySizeLod)
2706 (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_lod);
2707
2708 /* For OpFragmentFetchAMD, we always have a multisample index */
2709 if (opcode == SpvOpFragmentFetchAMD)
2710 (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_ms_index);
2711
2712 /* Now we need to handle some number of optional arguments */
2713 struct vtn_value *gather_offsets = NULL;
2714 if (idx < count) {
2715 uint32_t operands = w[idx];
2716
2717 if (operands & SpvImageOperandsBiasMask) {
2718 vtn_assert(texop == nir_texop_tex ||
2719 texop == nir_texop_tg4);
2720 if (texop == nir_texop_tex)
2721 texop = nir_texop_txb;
2722 uint32_t arg = image_operand_arg(b, w, count, idx,
2723 SpvImageOperandsBiasMask);
2724 (*p++) = vtn_tex_src(b, w[arg], nir_tex_src_bias);
2725 }
2726
2727 if (operands & SpvImageOperandsLodMask) {
2728 vtn_assert(texop == nir_texop_txl || texop == nir_texop_txf ||
2729 texop == nir_texop_txs || texop == nir_texop_tg4);
2730 uint32_t arg = image_operand_arg(b, w, count, idx,
2731 SpvImageOperandsLodMask);
2732 (*p++) = vtn_tex_src(b, w[arg], nir_tex_src_lod);
2733 }
2734
2735 if (operands & SpvImageOperandsGradMask) {
2736 vtn_assert(texop == nir_texop_txl);
2737 texop = nir_texop_txd;
2738 uint32_t arg = image_operand_arg(b, w, count, idx,
2739 SpvImageOperandsGradMask);
2740 (*p++) = vtn_tex_src(b, w[arg], nir_tex_src_ddx);
2741 (*p++) = vtn_tex_src(b, w[arg + 1], nir_tex_src_ddy);
2742 }
2743
2744 vtn_fail_if(util_bitcount(operands & (SpvImageOperandsConstOffsetsMask |
2745 SpvImageOperandsOffsetMask |
2746 SpvImageOperandsConstOffsetMask)) > 1,
2747 "At most one of the ConstOffset, Offset, and ConstOffsets "
2748 "image operands can be used on a given instruction.");
2749
2750 if (operands & SpvImageOperandsOffsetMask) {
2751 uint32_t arg = image_operand_arg(b, w, count, idx,
2752 SpvImageOperandsOffsetMask);
2753 (*p++) = vtn_tex_src(b, w[arg], nir_tex_src_offset);
2754 }
2755
2756 if (operands & SpvImageOperandsConstOffsetMask) {
2757 uint32_t arg = image_operand_arg(b, w, count, idx,
2758 SpvImageOperandsConstOffsetMask);
2759 (*p++) = vtn_tex_src(b, w[arg], nir_tex_src_offset);
2760 }
2761
2762 if (operands & SpvImageOperandsConstOffsetsMask) {
2763 vtn_assert(texop == nir_texop_tg4);
2764 uint32_t arg = image_operand_arg(b, w, count, idx,
2765 SpvImageOperandsConstOffsetsMask);
2766 gather_offsets = vtn_value(b, w[arg], vtn_value_type_constant);
2767 }
2768
2769 if (operands & SpvImageOperandsSampleMask) {
2770 vtn_assert(texop == nir_texop_txf_ms);
2771 uint32_t arg = image_operand_arg(b, w, count, idx,
2772 SpvImageOperandsSampleMask);
2773 texop = nir_texop_txf_ms;
2774 (*p++) = vtn_tex_src(b, w[arg], nir_tex_src_ms_index);
2775 }
2776
2777 if (operands & SpvImageOperandsMinLodMask) {
2778 vtn_assert(texop == nir_texop_tex ||
2779 texop == nir_texop_txb ||
2780 texop == nir_texop_txd);
2781 uint32_t arg = image_operand_arg(b, w, count, idx,
2782 SpvImageOperandsMinLodMask);
2783 (*p++) = vtn_tex_src(b, w[arg], nir_tex_src_min_lod);
2784 }
2785 }
2786
2787 nir_tex_instr *instr = nir_tex_instr_create(b->shader, p - srcs);
2788 instr->op = texop;
2789
2790 memcpy(instr->src, srcs, instr->num_srcs * sizeof(*instr->src));
2791
2792 instr->coord_components = coord_components;
2793 instr->sampler_dim = sampler_dim;
2794 instr->is_array = is_array;
2795 instr->is_shadow = is_shadow;
2796 instr->is_new_style_shadow =
2797 is_shadow && glsl_get_components(ret_type->type) == 1;
2798 instr->component = gather_component;
2799
2800 /* The Vulkan spec says:
2801 *
2802 * "If an instruction loads from or stores to a resource (including
2803 * atomics and image instructions) and the resource descriptor being
2804 * accessed is not dynamically uniform, then the operand corresponding
2805 * to that resource (e.g. the pointer or sampled image operand) must be
2806 * decorated with NonUniform."
2807 *
2808 * It's very careful to specify that the exact operand must be decorated
2809 * NonUniform. The SPIR-V parser is not expected to chase through long
2810 * chains to find the NonUniform decoration. It's either right there or we
2811 * can assume it doesn't exist.
2812 */
2813 enum gl_access_qualifier access = 0;
2814 vtn_foreach_decoration(b, sampled_val, non_uniform_decoration_cb, &access);
2815
2816 if (sampled_val->propagated_non_uniform)
2817 access |= ACCESS_NON_UNIFORM;
2818
2819 if (image && (access & ACCESS_NON_UNIFORM))
2820 instr->texture_non_uniform = true;
2821
2822 if (sampler && (access & ACCESS_NON_UNIFORM))
2823 instr->sampler_non_uniform = true;
2824
2825 /* for non-query ops, get dest_type from SPIR-V return type */
2826 if (dest_type == nir_type_invalid) {
2827 /* the return type should match the image type, unless the image type is
2828 * VOID (CL image), in which case the return type dictates the sampler
2829 */
2830 enum glsl_base_type sampler_base =
2831 glsl_get_sampler_result_type(image->type);
2832 enum glsl_base_type ret_base = glsl_get_base_type(ret_type->type);
2833 vtn_fail_if(sampler_base != ret_base && sampler_base != GLSL_TYPE_VOID,
2834 "SPIR-V return type mismatches image type. This is only valid "
2835 "for untyped images (OpenCL).");
2836 switch (ret_base) {
2837 case GLSL_TYPE_FLOAT: dest_type = nir_type_float; break;
2838 case GLSL_TYPE_INT: dest_type = nir_type_int; break;
2839 case GLSL_TYPE_UINT: dest_type = nir_type_uint; break;
2840 case GLSL_TYPE_BOOL: dest_type = nir_type_bool; break;
2841 default:
2842 vtn_fail("Invalid base type for sampler result");
2843 }
2844 }
2845
2846 instr->dest_type = dest_type;
2847
2848 nir_ssa_dest_init(&instr->instr, &instr->dest,
2849 nir_tex_instr_dest_size(instr), 32, NULL);
2850
2851 vtn_assert(glsl_get_vector_elements(ret_type->type) ==
2852 nir_tex_instr_dest_size(instr));
2853
2854 if (gather_offsets) {
2855 vtn_fail_if(gather_offsets->type->base_type != vtn_base_type_array ||
2856 gather_offsets->type->length != 4,
2857 "ConstOffsets must be an array of size four of vectors "
2858 "of two integer components");
2859
2860 struct vtn_type *vec_type = gather_offsets->type->array_element;
2861 vtn_fail_if(vec_type->base_type != vtn_base_type_vector ||
2862 vec_type->length != 2 ||
2863 !glsl_type_is_integer(vec_type->type),
2864 "ConstOffsets must be an array of size four of vectors "
2865 "of two integer components");
2866
2867 unsigned bit_size = glsl_get_bit_size(vec_type->type);
2868 for (uint32_t i = 0; i < 4; i++) {
2869 const nir_const_value *cvec =
2870 gather_offsets->constant->elements[i]->values;
2871 for (uint32_t j = 0; j < 2; j++) {
2872 switch (bit_size) {
2873 case 8: instr->tg4_offsets[i][j] = cvec[j].i8; break;
2874 case 16: instr->tg4_offsets[i][j] = cvec[j].i16; break;
2875 case 32: instr->tg4_offsets[i][j] = cvec[j].i32; break;
2876 case 64: instr->tg4_offsets[i][j] = cvec[j].i64; break;
2877 default:
2878 vtn_fail("Unsupported bit size: %u", bit_size);
2879 }
2880 }
2881 }
2882 }
2883
2884 nir_builder_instr_insert(&b->nb, &instr->instr);
2885
2886 vtn_push_nir_ssa(b, w[2], &instr->dest.ssa);
2887 }
2888
2889 static void
2890 fill_common_atomic_sources(struct vtn_builder *b, SpvOp opcode,
2891 const uint32_t *w, nir_src *src)
2892 {
2893 const struct glsl_type *type = vtn_get_type(b, w[1])->type;
2894 unsigned bit_size = glsl_get_bit_size(type);
2895
2896 switch (opcode) {
2897 case SpvOpAtomicIIncrement:
2898 src[0] = nir_src_for_ssa(nir_imm_intN_t(&b->nb, 1, bit_size));
2899 break;
2900
2901 case SpvOpAtomicIDecrement:
2902 src[0] = nir_src_for_ssa(nir_imm_intN_t(&b->nb, -1, bit_size));
2903 break;
2904
2905 case SpvOpAtomicISub:
2906 src[0] =
2907 nir_src_for_ssa(nir_ineg(&b->nb, vtn_get_nir_ssa(b, w[6])));
2908 break;
2909
2910 case SpvOpAtomicCompareExchange:
2911 case SpvOpAtomicCompareExchangeWeak:
2912 src[0] = nir_src_for_ssa(vtn_get_nir_ssa(b, w[8]));
2913 src[1] = nir_src_for_ssa(vtn_get_nir_ssa(b, w[7]));
2914 break;
2915
2916 case SpvOpAtomicExchange:
2917 case SpvOpAtomicIAdd:
2918 case SpvOpAtomicSMin:
2919 case SpvOpAtomicUMin:
2920 case SpvOpAtomicSMax:
2921 case SpvOpAtomicUMax:
2922 case SpvOpAtomicAnd:
2923 case SpvOpAtomicOr:
2924 case SpvOpAtomicXor:
2925 case SpvOpAtomicFAddEXT:
2926 src[0] = nir_src_for_ssa(vtn_get_nir_ssa(b, w[6]));
2927 break;
2928
2929 default:
2930 vtn_fail_with_opcode("Invalid SPIR-V atomic", opcode);
2931 }
2932 }
2933
2934 static nir_ssa_def *
2935 get_image_coord(struct vtn_builder *b, uint32_t value)
2936 {
2937 nir_ssa_def *coord = vtn_get_nir_ssa(b, value);
2938
2939 /* The image_load_store intrinsics assume a 4-dim coordinate */
2940 unsigned swizzle[4];
2941 for (unsigned i = 0; i < 4; i++)
2942 swizzle[i] = MIN2(i, coord->num_components - 1);
2943
2944 return nir_swizzle(&b->nb, coord, swizzle, 4);
2945 }
2946
2947 static nir_ssa_def *
2948 expand_to_vec4(nir_builder *b, nir_ssa_def *value)
2949 {
2950 if (value->num_components == 4)
2951 return value;
2952
2953 unsigned swiz[4];
2954 for (unsigned i = 0; i < 4; i++)
2955 swiz[i] = i < value->num_components ? i : 0;
2956 return nir_swizzle(b, value, swiz, 4);
2957 }
2958
2959 static void
2960 vtn_handle_image(struct vtn_builder *b, SpvOp opcode,
2961 const uint32_t *w, unsigned count)
2962 {
2963 /* Just get this one out of the way */
2964 if (opcode == SpvOpImageTexelPointer) {
2965 struct vtn_value *val =
2966 vtn_push_value(b, w[2], vtn_value_type_image_pointer);
2967 val->image = ralloc(b, struct vtn_image_pointer);
2968
2969 val->image->image = vtn_nir_deref(b, w[3]);
2970 val->image->coord = get_image_coord(b, w[4]);
2971 val->image->sample = vtn_get_nir_ssa(b, w[5]);
2972 val->image->lod = nir_imm_int(&b->nb, 0);
2973 return;
2974 }
2975
2976 struct vtn_image_pointer image;
2977 SpvScope scope = SpvScopeInvocation;
2978 SpvMemorySemanticsMask semantics = 0;
2979
2980 enum gl_access_qualifier access = 0;
2981
2982 struct vtn_value *res_val;
2983 switch (opcode) {
2984 case SpvOpAtomicExchange:
2985 case SpvOpAtomicCompareExchange:
2986 case SpvOpAtomicCompareExchangeWeak:
2987 case SpvOpAtomicIIncrement:
2988 case SpvOpAtomicIDecrement:
2989 case SpvOpAtomicIAdd:
2990 case SpvOpAtomicISub:
2991 case SpvOpAtomicLoad:
2992 case SpvOpAtomicSMin:
2993 case SpvOpAtomicUMin:
2994 case SpvOpAtomicSMax:
2995 case SpvOpAtomicUMax:
2996 case SpvOpAtomicAnd:
2997 case SpvOpAtomicOr:
2998 case SpvOpAtomicXor:
2999 case SpvOpAtomicFAddEXT:
3000 res_val = vtn_value(b, w[3], vtn_value_type_image_pointer);
3001 image = *res_val->image;
3002 scope = vtn_constant_uint(b, w[4]);
3003 semantics = vtn_constant_uint(b, w[5]);
3004 access |= ACCESS_COHERENT;
3005 break;
3006
3007 case SpvOpAtomicStore:
3008 res_val = vtn_value(b, w[1], vtn_value_type_image_pointer);
3009 image = *res_val->image;
3010 scope = vtn_constant_uint(b, w[2]);
3011 semantics = vtn_constant_uint(b, w[3]);
3012 access |= ACCESS_COHERENT;
3013 break;
3014
3015 case SpvOpImageQuerySizeLod:
3016 res_val = vtn_untyped_value(b, w[3]);
3017 image.image = vtn_get_image(b, w[3]);
3018 image.coord = NULL;
3019 image.sample = NULL;
3020 image.lod = vtn_ssa_value(b, w[4])->def;
3021 break;
3022
3023 case SpvOpImageQuerySize:
3024 res_val = vtn_untyped_value(b, w[3]);
3025 image.image = vtn_get_image(b, w[3]);
3026 image.coord = NULL;
3027 image.sample = NULL;
3028 image.lod = NULL;
3029 break;
3030
3031 case SpvOpImageQueryFormat:
3032 case SpvOpImageQueryOrder:
3033 res_val = vtn_untyped_value(b, w[3]);
3034 image.image = vtn_get_image(b, w[3]);
3035 image.coord = NULL;
3036 image.sample = NULL;
3037 image.lod = NULL;
3038 break;
3039
3040 case SpvOpImageRead: {
3041 res_val = vtn_untyped_value(b, w[3]);
3042 image.image = vtn_get_image(b, w[3]);
3043 image.coord = get_image_coord(b, w[4]);
3044
3045 const SpvImageOperandsMask operands =
3046 count > 5 ? w[5] : SpvImageOperandsMaskNone;
3047
3048 if (operands & SpvImageOperandsSampleMask) {
3049 uint32_t arg = image_operand_arg(b, w, count, 5,
3050 SpvImageOperandsSampleMask);
3051 image.sample = vtn_get_nir_ssa(b, w[arg]);
3052 } else {
3053 image.sample = nir_ssa_undef(&b->nb, 1, 32);
3054 }
3055
3056 if (operands & SpvImageOperandsMakeTexelVisibleMask) {
3057 vtn_fail_if((operands & SpvImageOperandsNonPrivateTexelMask) == 0,
3058 "MakeTexelVisible requires NonPrivateTexel to also be set.");
3059 uint32_t arg = image_operand_arg(b, w, count, 5,
3060 SpvImageOperandsMakeTexelVisibleMask);
3061 semantics = SpvMemorySemanticsMakeVisibleMask;
3062 scope = vtn_constant_uint(b, w[arg]);
3063 }
3064
3065 if (operands & SpvImageOperandsLodMask) {
3066 uint32_t arg = image_operand_arg(b, w, count, 5,
3067 SpvImageOperandsLodMask);
3068 image.lod = vtn_get_nir_ssa(b, w[arg]);
3069 } else {
3070 image.lod = nir_imm_int(&b->nb, 0);
3071 }
3072
3073 if (operands & SpvImageOperandsVolatileTexelMask)
3074 access |= ACCESS_VOLATILE;
3075
3076 break;
3077 }
3078
3079 case SpvOpImageWrite: {
3080 res_val = vtn_untyped_value(b, w[1]);
3081 image.image = vtn_get_image(b, w[1]);
3082 image.coord = get_image_coord(b, w[2]);
3083
3084 /* texel = w[3] */
3085
3086 const SpvImageOperandsMask operands =
3087 count > 4 ? w[4] : SpvImageOperandsMaskNone;
3088
3089 if (operands & SpvImageOperandsSampleMask) {
3090 uint32_t arg = image_operand_arg(b, w, count, 4,
3091 SpvImageOperandsSampleMask);
3092 image.sample = vtn_get_nir_ssa(b, w[arg]);
3093 } else {
3094 image.sample = nir_ssa_undef(&b->nb, 1, 32);
3095 }
3096
3097 if (operands & SpvImageOperandsMakeTexelAvailableMask) {
3098 vtn_fail_if((operands & SpvImageOperandsNonPrivateTexelMask) == 0,
3099 "MakeTexelAvailable requires NonPrivateTexel to also be set.");
3100 uint32_t arg = image_operand_arg(b, w, count, 4,
3101 SpvImageOperandsMakeTexelAvailableMask);
3102 semantics = SpvMemorySemanticsMakeAvailableMask;