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