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