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 "spirv_info.h"
33
34 struct spec_constant_value {
35 bool is_double;
36 union {
37 uint32_t data32;
38 uint64_t data64;
39 };
40 };
41
42 void
_vtn_warn(const char * file,int line,const char * msg,...)43 _vtn_warn(const char *file, int line, const char *msg, ...)
44 {
45 char *formatted;
46 va_list args;
47
48 va_start(args, msg);
49 formatted = ralloc_vasprintf(NULL, msg, args);
50 va_end(args);
51
52 fprintf(stderr, "%s:%d WARNING: %s\n", file, line, formatted);
53
54 ralloc_free(formatted);
55 }
56
57 static struct vtn_ssa_value *
vtn_undef_ssa_value(struct vtn_builder * b,const struct glsl_type * type)58 vtn_undef_ssa_value(struct vtn_builder *b, const struct glsl_type *type)
59 {
60 struct vtn_ssa_value *val = rzalloc(b, struct vtn_ssa_value);
61 val->type = type;
62
63 if (glsl_type_is_vector_or_scalar(type)) {
64 unsigned num_components = glsl_get_vector_elements(val->type);
65 unsigned bit_size = glsl_get_bit_size(val->type);
66 val->def = nir_ssa_undef(&b->nb, num_components, bit_size);
67 } else {
68 unsigned elems = glsl_get_length(val->type);
69 val->elems = ralloc_array(b, struct vtn_ssa_value *, elems);
70 if (glsl_type_is_matrix(type)) {
71 const struct glsl_type *elem_type =
72 glsl_vector_type(glsl_get_base_type(type),
73 glsl_get_vector_elements(type));
74
75 for (unsigned i = 0; i < elems; i++)
76 val->elems[i] = vtn_undef_ssa_value(b, elem_type);
77 } else if (glsl_type_is_array(type)) {
78 const struct glsl_type *elem_type = glsl_get_array_element(type);
79 for (unsigned i = 0; i < elems; i++)
80 val->elems[i] = vtn_undef_ssa_value(b, elem_type);
81 } else {
82 for (unsigned i = 0; i < elems; i++) {
83 const struct glsl_type *elem_type = glsl_get_struct_field(type, i);
84 val->elems[i] = vtn_undef_ssa_value(b, elem_type);
85 }
86 }
87 }
88
89 return val;
90 }
91
92 static struct vtn_ssa_value *
vtn_const_ssa_value(struct vtn_builder * b,nir_constant * constant,const struct glsl_type * type)93 vtn_const_ssa_value(struct vtn_builder *b, nir_constant *constant,
94 const struct glsl_type *type)
95 {
96 struct hash_entry *entry = _mesa_hash_table_search(b->const_table, constant);
97
98 if (entry)
99 return entry->data;
100
101 struct vtn_ssa_value *val = rzalloc(b, struct vtn_ssa_value);
102 val->type = type;
103
104 switch (glsl_get_base_type(type)) {
105 case GLSL_TYPE_INT:
106 case GLSL_TYPE_UINT:
107 case GLSL_TYPE_BOOL:
108 case GLSL_TYPE_FLOAT:
109 case GLSL_TYPE_DOUBLE: {
110 int bit_size = glsl_get_bit_size(type);
111 if (glsl_type_is_vector_or_scalar(type)) {
112 unsigned num_components = glsl_get_vector_elements(val->type);
113 nir_load_const_instr *load =
114 nir_load_const_instr_create(b->shader, num_components, bit_size);
115
116 load->value = constant->values[0];
117
118 nir_instr_insert_before_cf_list(&b->impl->body, &load->instr);
119 val->def = &load->def;
120 } else {
121 assert(glsl_type_is_matrix(type));
122 unsigned rows = glsl_get_vector_elements(val->type);
123 unsigned columns = glsl_get_matrix_columns(val->type);
124 val->elems = ralloc_array(b, struct vtn_ssa_value *, columns);
125
126 for (unsigned i = 0; i < columns; i++) {
127 struct vtn_ssa_value *col_val = rzalloc(b, struct vtn_ssa_value);
128 col_val->type = glsl_get_column_type(val->type);
129 nir_load_const_instr *load =
130 nir_load_const_instr_create(b->shader, rows, bit_size);
131
132 load->value = constant->values[i];
133
134 nir_instr_insert_before_cf_list(&b->impl->body, &load->instr);
135 col_val->def = &load->def;
136
137 val->elems[i] = col_val;
138 }
139 }
140 break;
141 }
142
143 case GLSL_TYPE_ARRAY: {
144 unsigned elems = glsl_get_length(val->type);
145 val->elems = ralloc_array(b, struct vtn_ssa_value *, elems);
146 const struct glsl_type *elem_type = glsl_get_array_element(val->type);
147 for (unsigned i = 0; i < elems; i++)
148 val->elems[i] = vtn_const_ssa_value(b, constant->elements[i],
149 elem_type);
150 break;
151 }
152
153 case GLSL_TYPE_STRUCT: {
154 unsigned elems = glsl_get_length(val->type);
155 val->elems = ralloc_array(b, struct vtn_ssa_value *, elems);
156 for (unsigned i = 0; i < elems; i++) {
157 const struct glsl_type *elem_type =
158 glsl_get_struct_field(val->type, i);
159 val->elems[i] = vtn_const_ssa_value(b, constant->elements[i],
160 elem_type);
161 }
162 break;
163 }
164
165 default:
166 unreachable("bad constant type");
167 }
168
169 return val;
170 }
171
172 struct vtn_ssa_value *
vtn_ssa_value(struct vtn_builder * b,uint32_t value_id)173 vtn_ssa_value(struct vtn_builder *b, uint32_t value_id)
174 {
175 struct vtn_value *val = vtn_untyped_value(b, value_id);
176 switch (val->value_type) {
177 case vtn_value_type_undef:
178 return vtn_undef_ssa_value(b, val->type->type);
179
180 case vtn_value_type_constant:
181 return vtn_const_ssa_value(b, val->constant, val->const_type);
182
183 case vtn_value_type_ssa:
184 return val->ssa;
185
186 case vtn_value_type_access_chain:
187 /* This is needed for function parameters */
188 return vtn_variable_load(b, val->access_chain);
189
190 default:
191 unreachable("Invalid type for an SSA value");
192 }
193 }
194
195 static char *
vtn_string_literal(struct vtn_builder * b,const uint32_t * words,unsigned word_count,unsigned * words_used)196 vtn_string_literal(struct vtn_builder *b, const uint32_t *words,
197 unsigned word_count, unsigned *words_used)
198 {
199 char *dup = ralloc_strndup(b, (char *)words, word_count * sizeof(*words));
200 if (words_used) {
201 /* Ammount of space taken by the string (including the null) */
202 unsigned len = strlen(dup) + 1;
203 *words_used = DIV_ROUND_UP(len, sizeof(*words));
204 }
205 return dup;
206 }
207
208 const uint32_t *
vtn_foreach_instruction(struct vtn_builder * b,const uint32_t * start,const uint32_t * end,vtn_instruction_handler handler)209 vtn_foreach_instruction(struct vtn_builder *b, const uint32_t *start,
210 const uint32_t *end, vtn_instruction_handler handler)
211 {
212 b->file = NULL;
213 b->line = -1;
214 b->col = -1;
215
216 const uint32_t *w = start;
217 while (w < end) {
218 SpvOp opcode = w[0] & SpvOpCodeMask;
219 unsigned count = w[0] >> SpvWordCountShift;
220 assert(count >= 1 && w + count <= end);
221
222 switch (opcode) {
223 case SpvOpNop:
224 break; /* Do nothing */
225
226 case SpvOpLine:
227 b->file = vtn_value(b, w[1], vtn_value_type_string)->str;
228 b->line = w[2];
229 b->col = w[3];
230 break;
231
232 case SpvOpNoLine:
233 b->file = NULL;
234 b->line = -1;
235 b->col = -1;
236 break;
237
238 default:
239 if (!handler(b, opcode, w, count))
240 return w;
241 break;
242 }
243
244 w += count;
245 }
246 assert(w == end);
247 return w;
248 }
249
250 static void
vtn_handle_extension(struct vtn_builder * b,SpvOp opcode,const uint32_t * w,unsigned count)251 vtn_handle_extension(struct vtn_builder *b, SpvOp opcode,
252 const uint32_t *w, unsigned count)
253 {
254 switch (opcode) {
255 case SpvOpExtInstImport: {
256 struct vtn_value *val = vtn_push_value(b, w[1], vtn_value_type_extension);
257 if (strcmp((const char *)&w[2], "GLSL.std.450") == 0) {
258 val->ext_handler = vtn_handle_glsl450_instruction;
259 } else {
260 assert(!"Unsupported extension");
261 }
262 break;
263 }
264
265 case SpvOpExtInst: {
266 struct vtn_value *val = vtn_value(b, w[3], vtn_value_type_extension);
267 bool handled = val->ext_handler(b, w[4], w, count);
268 (void)handled;
269 assert(handled);
270 break;
271 }
272
273 default:
274 unreachable("Unhandled opcode");
275 }
276 }
277
278 static void
_foreach_decoration_helper(struct vtn_builder * b,struct vtn_value * base_value,int parent_member,struct vtn_value * value,vtn_decoration_foreach_cb cb,void * data)279 _foreach_decoration_helper(struct vtn_builder *b,
280 struct vtn_value *base_value,
281 int parent_member,
282 struct vtn_value *value,
283 vtn_decoration_foreach_cb cb, void *data)
284 {
285 for (struct vtn_decoration *dec = value->decoration; dec; dec = dec->next) {
286 int member;
287 if (dec->scope == VTN_DEC_DECORATION) {
288 member = parent_member;
289 } else if (dec->scope >= VTN_DEC_STRUCT_MEMBER0) {
290 assert(parent_member == -1);
291 member = dec->scope - VTN_DEC_STRUCT_MEMBER0;
292 } else {
293 /* Not a decoration */
294 continue;
295 }
296
297 if (dec->group) {
298 assert(dec->group->value_type == vtn_value_type_decoration_group);
299 _foreach_decoration_helper(b, base_value, member, dec->group,
300 cb, data);
301 } else {
302 cb(b, base_value, member, dec, data);
303 }
304 }
305 }
306
307 /** Iterates (recursively if needed) over all of the decorations on a value
308 *
309 * This function iterates over all of the decorations applied to a given
310 * value. If it encounters a decoration group, it recurses into the group
311 * and iterates over all of those decorations as well.
312 */
313 void
vtn_foreach_decoration(struct vtn_builder * b,struct vtn_value * value,vtn_decoration_foreach_cb cb,void * data)314 vtn_foreach_decoration(struct vtn_builder *b, struct vtn_value *value,
315 vtn_decoration_foreach_cb cb, void *data)
316 {
317 _foreach_decoration_helper(b, value, -1, value, cb, data);
318 }
319
320 void
vtn_foreach_execution_mode(struct vtn_builder * b,struct vtn_value * value,vtn_execution_mode_foreach_cb cb,void * data)321 vtn_foreach_execution_mode(struct vtn_builder *b, struct vtn_value *value,
322 vtn_execution_mode_foreach_cb cb, void *data)
323 {
324 for (struct vtn_decoration *dec = value->decoration; dec; dec = dec->next) {
325 if (dec->scope != VTN_DEC_EXECUTION_MODE)
326 continue;
327
328 assert(dec->group == NULL);
329 cb(b, value, dec, data);
330 }
331 }
332
333 static void
vtn_handle_decoration(struct vtn_builder * b,SpvOp opcode,const uint32_t * w,unsigned count)334 vtn_handle_decoration(struct vtn_builder *b, SpvOp opcode,
335 const uint32_t *w, unsigned count)
336 {
337 const uint32_t *w_end = w + count;
338 const uint32_t target = w[1];
339 w += 2;
340
341 switch (opcode) {
342 case SpvOpDecorationGroup:
343 vtn_push_value(b, target, vtn_value_type_decoration_group);
344 break;
345
346 case SpvOpDecorate:
347 case SpvOpMemberDecorate:
348 case SpvOpExecutionMode: {
349 struct vtn_value *val = &b->values[target];
350
351 struct vtn_decoration *dec = rzalloc(b, struct vtn_decoration);
352 switch (opcode) {
353 case SpvOpDecorate:
354 dec->scope = VTN_DEC_DECORATION;
355 break;
356 case SpvOpMemberDecorate:
357 dec->scope = VTN_DEC_STRUCT_MEMBER0 + *(w++);
358 break;
359 case SpvOpExecutionMode:
360 dec->scope = VTN_DEC_EXECUTION_MODE;
361 break;
362 default:
363 unreachable("Invalid decoration opcode");
364 }
365 dec->decoration = *(w++);
366 dec->literals = w;
367
368 /* Link into the list */
369 dec->next = val->decoration;
370 val->decoration = dec;
371 break;
372 }
373
374 case SpvOpGroupMemberDecorate:
375 case SpvOpGroupDecorate: {
376 struct vtn_value *group =
377 vtn_value(b, target, vtn_value_type_decoration_group);
378
379 for (; w < w_end; w++) {
380 struct vtn_value *val = vtn_untyped_value(b, *w);
381 struct vtn_decoration *dec = rzalloc(b, struct vtn_decoration);
382
383 dec->group = group;
384 if (opcode == SpvOpGroupDecorate) {
385 dec->scope = VTN_DEC_DECORATION;
386 } else {
387 dec->scope = VTN_DEC_STRUCT_MEMBER0 + *(++w);
388 }
389
390 /* Link into the list */
391 dec->next = val->decoration;
392 val->decoration = dec;
393 }
394 break;
395 }
396
397 default:
398 unreachable("Unhandled opcode");
399 }
400 }
401
402 struct member_decoration_ctx {
403 unsigned num_fields;
404 struct glsl_struct_field *fields;
405 struct vtn_type *type;
406 };
407
408 /* does a shallow copy of a vtn_type */
409
410 static struct vtn_type *
vtn_type_copy(struct vtn_builder * b,struct vtn_type * src)411 vtn_type_copy(struct vtn_builder *b, struct vtn_type *src)
412 {
413 struct vtn_type *dest = ralloc(b, struct vtn_type);
414 dest->type = src->type;
415 dest->is_builtin = src->is_builtin;
416 if (src->is_builtin)
417 dest->builtin = src->builtin;
418
419 if (!glsl_type_is_scalar(src->type)) {
420 switch (glsl_get_base_type(src->type)) {
421 case GLSL_TYPE_INT:
422 case GLSL_TYPE_UINT:
423 case GLSL_TYPE_BOOL:
424 case GLSL_TYPE_FLOAT:
425 case GLSL_TYPE_DOUBLE:
426 case GLSL_TYPE_ARRAY:
427 dest->row_major = src->row_major;
428 dest->stride = src->stride;
429 dest->array_element = src->array_element;
430 break;
431
432 case GLSL_TYPE_STRUCT: {
433 unsigned elems = glsl_get_length(src->type);
434
435 dest->members = ralloc_array(b, struct vtn_type *, elems);
436 memcpy(dest->members, src->members, elems * sizeof(struct vtn_type *));
437
438 dest->offsets = ralloc_array(b, unsigned, elems);
439 memcpy(dest->offsets, src->offsets, elems * sizeof(unsigned));
440 break;
441 }
442
443 default:
444 unreachable("unhandled type");
445 }
446 }
447
448 return dest;
449 }
450
451 static struct vtn_type *
mutable_matrix_member(struct vtn_builder * b,struct vtn_type * type,int member)452 mutable_matrix_member(struct vtn_builder *b, struct vtn_type *type, int member)
453 {
454 type->members[member] = vtn_type_copy(b, type->members[member]);
455 type = type->members[member];
456
457 /* We may have an array of matrices.... Oh, joy! */
458 while (glsl_type_is_array(type->type)) {
459 type->array_element = vtn_type_copy(b, type->array_element);
460 type = type->array_element;
461 }
462
463 assert(glsl_type_is_matrix(type->type));
464
465 return type;
466 }
467
468 static void
struct_member_decoration_cb(struct vtn_builder * b,struct vtn_value * val,int member,const struct vtn_decoration * dec,void * void_ctx)469 struct_member_decoration_cb(struct vtn_builder *b,
470 struct vtn_value *val, int member,
471 const struct vtn_decoration *dec, void *void_ctx)
472 {
473 struct member_decoration_ctx *ctx = void_ctx;
474
475 if (member < 0)
476 return;
477
478 assert(member < ctx->num_fields);
479
480 switch (dec->decoration) {
481 case SpvDecorationNonWritable:
482 case SpvDecorationNonReadable:
483 case SpvDecorationRelaxedPrecision:
484 case SpvDecorationVolatile:
485 case SpvDecorationCoherent:
486 case SpvDecorationUniform:
487 break; /* FIXME: Do nothing with this for now. */
488 case SpvDecorationNoPerspective:
489 ctx->fields[member].interpolation = INTERP_MODE_NOPERSPECTIVE;
490 break;
491 case SpvDecorationFlat:
492 ctx->fields[member].interpolation = INTERP_MODE_FLAT;
493 break;
494 case SpvDecorationCentroid:
495 ctx->fields[member].centroid = true;
496 break;
497 case SpvDecorationSample:
498 ctx->fields[member].sample = true;
499 break;
500 case SpvDecorationStream:
501 /* Vulkan only allows one GS stream */
502 assert(dec->literals[0] == 0);
503 break;
504 case SpvDecorationLocation:
505 ctx->fields[member].location = dec->literals[0];
506 break;
507 case SpvDecorationComponent:
508 break; /* FIXME: What should we do with these? */
509 case SpvDecorationBuiltIn:
510 ctx->type->members[member] = vtn_type_copy(b, ctx->type->members[member]);
511 ctx->type->members[member]->is_builtin = true;
512 ctx->type->members[member]->builtin = dec->literals[0];
513 ctx->type->builtin_block = true;
514 break;
515 case SpvDecorationOffset:
516 ctx->type->offsets[member] = dec->literals[0];
517 break;
518 case SpvDecorationMatrixStride:
519 mutable_matrix_member(b, ctx->type, member)->stride = dec->literals[0];
520 break;
521 case SpvDecorationColMajor:
522 break; /* Nothing to do here. Column-major is the default. */
523 case SpvDecorationRowMajor:
524 mutable_matrix_member(b, ctx->type, member)->row_major = true;
525 break;
526
527 case SpvDecorationPatch:
528 break;
529
530 case SpvDecorationSpecId:
531 case SpvDecorationBlock:
532 case SpvDecorationBufferBlock:
533 case SpvDecorationArrayStride:
534 case SpvDecorationGLSLShared:
535 case SpvDecorationGLSLPacked:
536 case SpvDecorationInvariant:
537 case SpvDecorationRestrict:
538 case SpvDecorationAliased:
539 case SpvDecorationConstant:
540 case SpvDecorationIndex:
541 case SpvDecorationBinding:
542 case SpvDecorationDescriptorSet:
543 case SpvDecorationLinkageAttributes:
544 case SpvDecorationNoContraction:
545 case SpvDecorationInputAttachmentIndex:
546 vtn_warn("Decoration not allowed on struct members: %s",
547 spirv_decoration_to_string(dec->decoration));
548 break;
549
550 case SpvDecorationXfbBuffer:
551 case SpvDecorationXfbStride:
552 vtn_warn("Vulkan does not have transform feedback");
553 break;
554
555 case SpvDecorationCPacked:
556 case SpvDecorationSaturatedConversion:
557 case SpvDecorationFuncParamAttr:
558 case SpvDecorationFPRoundingMode:
559 case SpvDecorationFPFastMathMode:
560 case SpvDecorationAlignment:
561 vtn_warn("Decoration only allowed for CL-style kernels: %s",
562 spirv_decoration_to_string(dec->decoration));
563 break;
564 }
565 }
566
567 static void
type_decoration_cb(struct vtn_builder * b,struct vtn_value * val,int member,const struct vtn_decoration * dec,void * ctx)568 type_decoration_cb(struct vtn_builder *b,
569 struct vtn_value *val, int member,
570 const struct vtn_decoration *dec, void *ctx)
571 {
572 struct vtn_type *type = val->type;
573
574 if (member != -1)
575 return;
576
577 switch (dec->decoration) {
578 case SpvDecorationArrayStride:
579 type->stride = dec->literals[0];
580 break;
581 case SpvDecorationBlock:
582 type->block = true;
583 break;
584 case SpvDecorationBufferBlock:
585 type->buffer_block = true;
586 break;
587 case SpvDecorationGLSLShared:
588 case SpvDecorationGLSLPacked:
589 /* Ignore these, since we get explicit offsets anyways */
590 break;
591
592 case SpvDecorationRowMajor:
593 case SpvDecorationColMajor:
594 case SpvDecorationMatrixStride:
595 case SpvDecorationBuiltIn:
596 case SpvDecorationNoPerspective:
597 case SpvDecorationFlat:
598 case SpvDecorationPatch:
599 case SpvDecorationCentroid:
600 case SpvDecorationSample:
601 case SpvDecorationVolatile:
602 case SpvDecorationCoherent:
603 case SpvDecorationNonWritable:
604 case SpvDecorationNonReadable:
605 case SpvDecorationUniform:
606 case SpvDecorationStream:
607 case SpvDecorationLocation:
608 case SpvDecorationComponent:
609 case SpvDecorationOffset:
610 case SpvDecorationXfbBuffer:
611 case SpvDecorationXfbStride:
612 vtn_warn("Decoraiton only allowed for struct members: %s",
613 spirv_decoration_to_string(dec->decoration));
614 break;
615
616 case SpvDecorationRelaxedPrecision:
617 case SpvDecorationSpecId:
618 case SpvDecorationInvariant:
619 case SpvDecorationRestrict:
620 case SpvDecorationAliased:
621 case SpvDecorationConstant:
622 case SpvDecorationIndex:
623 case SpvDecorationBinding:
624 case SpvDecorationDescriptorSet:
625 case SpvDecorationLinkageAttributes:
626 case SpvDecorationNoContraction:
627 case SpvDecorationInputAttachmentIndex:
628 vtn_warn("Decoraiton not allowed on types: %s",
629 spirv_decoration_to_string(dec->decoration));
630 break;
631
632 case SpvDecorationCPacked:
633 case SpvDecorationSaturatedConversion:
634 case SpvDecorationFuncParamAttr:
635 case SpvDecorationFPRoundingMode:
636 case SpvDecorationFPFastMathMode:
637 case SpvDecorationAlignment:
638 vtn_warn("Decoraiton only allowed for CL-style kernels: %s",
639 spirv_decoration_to_string(dec->decoration));
640 break;
641 }
642 }
643
644 static unsigned
translate_image_format(SpvImageFormat format)645 translate_image_format(SpvImageFormat format)
646 {
647 switch (format) {
648 case SpvImageFormatUnknown: return 0; /* GL_NONE */
649 case SpvImageFormatRgba32f: return 0x8814; /* GL_RGBA32F */
650 case SpvImageFormatRgba16f: return 0x881A; /* GL_RGBA16F */
651 case SpvImageFormatR32f: return 0x822E; /* GL_R32F */
652 case SpvImageFormatRgba8: return 0x8058; /* GL_RGBA8 */
653 case SpvImageFormatRgba8Snorm: return 0x8F97; /* GL_RGBA8_SNORM */
654 case SpvImageFormatRg32f: return 0x8230; /* GL_RG32F */
655 case SpvImageFormatRg16f: return 0x822F; /* GL_RG16F */
656 case SpvImageFormatR11fG11fB10f: return 0x8C3A; /* GL_R11F_G11F_B10F */
657 case SpvImageFormatR16f: return 0x822D; /* GL_R16F */
658 case SpvImageFormatRgba16: return 0x805B; /* GL_RGBA16 */
659 case SpvImageFormatRgb10A2: return 0x8059; /* GL_RGB10_A2 */
660 case SpvImageFormatRg16: return 0x822C; /* GL_RG16 */
661 case SpvImageFormatRg8: return 0x822B; /* GL_RG8 */
662 case SpvImageFormatR16: return 0x822A; /* GL_R16 */
663 case SpvImageFormatR8: return 0x8229; /* GL_R8 */
664 case SpvImageFormatRgba16Snorm: return 0x8F9B; /* GL_RGBA16_SNORM */
665 case SpvImageFormatRg16Snorm: return 0x8F99; /* GL_RG16_SNORM */
666 case SpvImageFormatRg8Snorm: return 0x8F95; /* GL_RG8_SNORM */
667 case SpvImageFormatR16Snorm: return 0x8F98; /* GL_R16_SNORM */
668 case SpvImageFormatR8Snorm: return 0x8F94; /* GL_R8_SNORM */
669 case SpvImageFormatRgba32i: return 0x8D82; /* GL_RGBA32I */
670 case SpvImageFormatRgba16i: return 0x8D88; /* GL_RGBA16I */
671 case SpvImageFormatRgba8i: return 0x8D8E; /* GL_RGBA8I */
672 case SpvImageFormatR32i: return 0x8235; /* GL_R32I */
673 case SpvImageFormatRg32i: return 0x823B; /* GL_RG32I */
674 case SpvImageFormatRg16i: return 0x8239; /* GL_RG16I */
675 case SpvImageFormatRg8i: return 0x8237; /* GL_RG8I */
676 case SpvImageFormatR16i: return 0x8233; /* GL_R16I */
677 case SpvImageFormatR8i: return 0x8231; /* GL_R8I */
678 case SpvImageFormatRgba32ui: return 0x8D70; /* GL_RGBA32UI */
679 case SpvImageFormatRgba16ui: return 0x8D76; /* GL_RGBA16UI */
680 case SpvImageFormatRgba8ui: return 0x8D7C; /* GL_RGBA8UI */
681 case SpvImageFormatR32ui: return 0x8236; /* GL_R32UI */
682 case SpvImageFormatRgb10a2ui: return 0x906F; /* GL_RGB10_A2UI */
683 case SpvImageFormatRg32ui: return 0x823C; /* GL_RG32UI */
684 case SpvImageFormatRg16ui: return 0x823A; /* GL_RG16UI */
685 case SpvImageFormatRg8ui: return 0x8238; /* GL_RG8UI */
686 case SpvImageFormatR16ui: return 0x823A; /* GL_RG16UI */
687 case SpvImageFormatR8ui: return 0x8232; /* GL_R8UI */
688 default:
689 assert(!"Invalid image format");
690 return 0;
691 }
692 }
693
694 static void
vtn_handle_type(struct vtn_builder * b,SpvOp opcode,const uint32_t * w,unsigned count)695 vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
696 const uint32_t *w, unsigned count)
697 {
698 struct vtn_value *val = vtn_push_value(b, w[1], vtn_value_type_type);
699
700 val->type = rzalloc(b, struct vtn_type);
701 val->type->is_builtin = false;
702 val->type->val = val;
703
704 switch (opcode) {
705 case SpvOpTypeVoid:
706 val->type->type = glsl_void_type();
707 break;
708 case SpvOpTypeBool:
709 val->type->type = glsl_bool_type();
710 break;
711 case SpvOpTypeInt: {
712 const bool signedness = w[3];
713 val->type->type = (signedness ? glsl_int_type() : glsl_uint_type());
714 break;
715 }
716 case SpvOpTypeFloat: {
717 int bit_size = w[2];
718 val->type->type = bit_size == 64 ? glsl_double_type() : glsl_float_type();
719 break;
720 }
721
722 case SpvOpTypeVector: {
723 struct vtn_type *base = vtn_value(b, w[2], vtn_value_type_type)->type;
724 unsigned elems = w[3];
725
726 assert(glsl_type_is_scalar(base->type));
727 val->type->type = glsl_vector_type(glsl_get_base_type(base->type), elems);
728
729 /* Vectors implicitly have sizeof(base_type) stride. For now, this
730 * is always 4 bytes. This will have to change if we want to start
731 * supporting doubles or half-floats.
732 */
733 val->type->stride = 4;
734 val->type->array_element = base;
735 break;
736 }
737
738 case SpvOpTypeMatrix: {
739 struct vtn_type *base = vtn_value(b, w[2], vtn_value_type_type)->type;
740 unsigned columns = w[3];
741
742 assert(glsl_type_is_vector(base->type));
743 val->type->type = glsl_matrix_type(glsl_get_base_type(base->type),
744 glsl_get_vector_elements(base->type),
745 columns);
746 assert(!glsl_type_is_error(val->type->type));
747 val->type->array_element = base;
748 val->type->row_major = false;
749 val->type->stride = 0;
750 break;
751 }
752
753 case SpvOpTypeRuntimeArray:
754 case SpvOpTypeArray: {
755 struct vtn_type *array_element =
756 vtn_value(b, w[2], vtn_value_type_type)->type;
757
758 unsigned length;
759 if (opcode == SpvOpTypeRuntimeArray) {
760 /* A length of 0 is used to denote unsized arrays */
761 length = 0;
762 } else {
763 length =
764 vtn_value(b, w[3], vtn_value_type_constant)->constant->values[0].u32[0];
765 }
766
767 val->type->type = glsl_array_type(array_element->type, length);
768 val->type->array_element = array_element;
769 val->type->stride = 0;
770 break;
771 }
772
773 case SpvOpTypeStruct: {
774 unsigned num_fields = count - 2;
775 val->type->members = ralloc_array(b, struct vtn_type *, num_fields);
776 val->type->offsets = ralloc_array(b, unsigned, num_fields);
777
778 NIR_VLA(struct glsl_struct_field, fields, count);
779 for (unsigned i = 0; i < num_fields; i++) {
780 val->type->members[i] =
781 vtn_value(b, w[i + 2], vtn_value_type_type)->type;
782 fields[i] = (struct glsl_struct_field) {
783 .type = val->type->members[i]->type,
784 .name = ralloc_asprintf(b, "field%d", i),
785 .location = -1,
786 };
787 }
788
789 struct member_decoration_ctx ctx = {
790 .num_fields = num_fields,
791 .fields = fields,
792 .type = val->type
793 };
794
795 vtn_foreach_decoration(b, val, struct_member_decoration_cb, &ctx);
796
797 const char *name = val->name ? val->name : "struct";
798
799 val->type->type = glsl_struct_type(fields, num_fields, name);
800 break;
801 }
802
803 case SpvOpTypeFunction: {
804 const struct glsl_type *return_type =
805 vtn_value(b, w[2], vtn_value_type_type)->type->type;
806 NIR_VLA(struct glsl_function_param, params, count - 3);
807 for (unsigned i = 0; i < count - 3; i++) {
808 params[i].type = vtn_value(b, w[i + 3], vtn_value_type_type)->type->type;
809
810 /* FIXME: */
811 params[i].in = true;
812 params[i].out = true;
813 }
814 val->type->type = glsl_function_type(return_type, params, count - 3);
815 break;
816 }
817
818 case SpvOpTypePointer:
819 /* FIXME: For now, we'll just do the really lame thing and return
820 * the same type. The validator should ensure that the proper number
821 * of dereferences happen
822 */
823 val->type = vtn_value(b, w[3], vtn_value_type_type)->type;
824 break;
825
826 case SpvOpTypeImage: {
827 const struct glsl_type *sampled_type =
828 vtn_value(b, w[2], vtn_value_type_type)->type->type;
829
830 assert(glsl_type_is_vector_or_scalar(sampled_type));
831
832 enum glsl_sampler_dim dim;
833 switch ((SpvDim)w[3]) {
834 case SpvDim1D: dim = GLSL_SAMPLER_DIM_1D; break;
835 case SpvDim2D: dim = GLSL_SAMPLER_DIM_2D; break;
836 case SpvDim3D: dim = GLSL_SAMPLER_DIM_3D; break;
837 case SpvDimCube: dim = GLSL_SAMPLER_DIM_CUBE; break;
838 case SpvDimRect: dim = GLSL_SAMPLER_DIM_RECT; break;
839 case SpvDimBuffer: dim = GLSL_SAMPLER_DIM_BUF; break;
840 case SpvDimSubpassData: dim = GLSL_SAMPLER_DIM_SUBPASS; break;
841 default:
842 unreachable("Invalid SPIR-V Sampler dimension");
843 }
844
845 bool is_shadow = w[4];
846 bool is_array = w[5];
847 bool multisampled = w[6];
848 unsigned sampled = w[7];
849 SpvImageFormat format = w[8];
850
851 if (count > 9)
852 val->type->access_qualifier = w[9];
853 else
854 val->type->access_qualifier = SpvAccessQualifierReadWrite;
855
856 if (multisampled) {
857 assert(dim == GLSL_SAMPLER_DIM_2D);
858 dim = GLSL_SAMPLER_DIM_MS;
859 }
860
861 val->type->image_format = translate_image_format(format);
862
863 if (sampled == 1) {
864 val->type->type = glsl_sampler_type(dim, is_shadow, is_array,
865 glsl_get_base_type(sampled_type));
866 } else if (sampled == 2) {
867 assert((dim == GLSL_SAMPLER_DIM_SUBPASS) || format);
868 assert(!is_shadow);
869 val->type->type = glsl_image_type(dim, is_array,
870 glsl_get_base_type(sampled_type));
871 } else {
872 assert(!"We need to know if the image will be sampled");
873 }
874 break;
875 }
876
877 case SpvOpTypeSampledImage:
878 val->type = vtn_value(b, w[2], vtn_value_type_type)->type;
879 break;
880
881 case SpvOpTypeSampler:
882 /* The actual sampler type here doesn't really matter. It gets
883 * thrown away the moment you combine it with an image. What really
884 * matters is that it's a sampler type as opposed to an integer type
885 * so the backend knows what to do.
886 */
887 val->type->type = glsl_bare_sampler_type();
888 break;
889
890 case SpvOpTypeOpaque:
891 case SpvOpTypeEvent:
892 case SpvOpTypeDeviceEvent:
893 case SpvOpTypeReserveId:
894 case SpvOpTypeQueue:
895 case SpvOpTypePipe:
896 default:
897 unreachable("Unhandled opcode");
898 }
899
900 vtn_foreach_decoration(b, val, type_decoration_cb, NULL);
901 }
902
903 static nir_constant *
vtn_null_constant(struct vtn_builder * b,const struct glsl_type * type)904 vtn_null_constant(struct vtn_builder *b, const struct glsl_type *type)
905 {
906 nir_constant *c = rzalloc(b, nir_constant);
907
908 switch (glsl_get_base_type(type)) {
909 case GLSL_TYPE_INT:
910 case GLSL_TYPE_UINT:
911 case GLSL_TYPE_BOOL:
912 case GLSL_TYPE_FLOAT:
913 case GLSL_TYPE_DOUBLE:
914 /* Nothing to do here. It's already initialized to zero */
915 break;
916
917 case GLSL_TYPE_ARRAY:
918 assert(glsl_get_length(type) > 0);
919 c->num_elements = glsl_get_length(type);
920 c->elements = ralloc_array(b, nir_constant *, c->num_elements);
921
922 c->elements[0] = vtn_null_constant(b, glsl_get_array_element(type));
923 for (unsigned i = 1; i < c->num_elements; i++)
924 c->elements[i] = c->elements[0];
925 break;
926
927 case GLSL_TYPE_STRUCT:
928 c->num_elements = glsl_get_length(type);
929 c->elements = ralloc_array(b, nir_constant *, c->num_elements);
930
931 for (unsigned i = 0; i < c->num_elements; i++) {
932 c->elements[i] = vtn_null_constant(b, glsl_get_struct_field(type, i));
933 }
934 break;
935
936 default:
937 unreachable("Invalid type for null constant");
938 }
939
940 return c;
941 }
942
943 static void
spec_constant_decoration_cb(struct vtn_builder * b,struct vtn_value * v,int member,const struct vtn_decoration * dec,void * data)944 spec_constant_decoration_cb(struct vtn_builder *b, struct vtn_value *v,
945 int member, const struct vtn_decoration *dec,
946 void *data)
947 {
948 assert(member == -1);
949 if (dec->decoration != SpvDecorationSpecId)
950 return;
951
952 struct spec_constant_value *const_value = data;
953
954 for (unsigned i = 0; i < b->num_specializations; i++) {
955 if (b->specializations[i].id == dec->literals[0]) {
956 if (const_value->is_double)
957 const_value->data64 = b->specializations[i].data64;
958 else
959 const_value->data32 = b->specializations[i].data32;
960 return;
961 }
962 }
963 }
964
965 static uint32_t
get_specialization(struct vtn_builder * b,struct vtn_value * val,uint32_t const_value)966 get_specialization(struct vtn_builder *b, struct vtn_value *val,
967 uint32_t const_value)
968 {
969 struct spec_constant_value data;
970 data.is_double = false;
971 data.data32 = const_value;
972 vtn_foreach_decoration(b, val, spec_constant_decoration_cb, &data);
973 return data.data32;
974 }
975
976 static uint64_t
get_specialization64(struct vtn_builder * b,struct vtn_value * val,uint64_t const_value)977 get_specialization64(struct vtn_builder *b, struct vtn_value *val,
978 uint64_t const_value)
979 {
980 struct spec_constant_value data;
981 data.is_double = true;
982 data.data64 = const_value;
983 vtn_foreach_decoration(b, val, spec_constant_decoration_cb, &data);
984 return data.data64;
985 }
986
987 static void
handle_workgroup_size_decoration_cb(struct vtn_builder * b,struct vtn_value * val,int member,const struct vtn_decoration * dec,void * data)988 handle_workgroup_size_decoration_cb(struct vtn_builder *b,
989 struct vtn_value *val,
990 int member,
991 const struct vtn_decoration *dec,
992 void *data)
993 {
994 assert(member == -1);
995 if (dec->decoration != SpvDecorationBuiltIn ||
996 dec->literals[0] != SpvBuiltInWorkgroupSize)
997 return;
998
999 assert(val->const_type == glsl_vector_type(GLSL_TYPE_UINT, 3));
1000
1001 b->shader->info->cs.local_size[0] = val->constant->values[0].u32[0];
1002 b->shader->info->cs.local_size[1] = val->constant->values[0].u32[1];
1003 b->shader->info->cs.local_size[2] = val->constant->values[0].u32[2];
1004 }
1005
1006 static void
vtn_handle_constant(struct vtn_builder * b,SpvOp opcode,const uint32_t * w,unsigned count)1007 vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
1008 const uint32_t *w, unsigned count)
1009 {
1010 struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_constant);
1011 val->const_type = vtn_value(b, w[1], vtn_value_type_type)->type->type;
1012 val->constant = rzalloc(b, nir_constant);
1013 switch (opcode) {
1014 case SpvOpConstantTrue:
1015 assert(val->const_type == glsl_bool_type());
1016 val->constant->values[0].u32[0] = NIR_TRUE;
1017 break;
1018 case SpvOpConstantFalse:
1019 assert(val->const_type == glsl_bool_type());
1020 val->constant->values[0].u32[0] = NIR_FALSE;
1021 break;
1022
1023 case SpvOpSpecConstantTrue:
1024 case SpvOpSpecConstantFalse: {
1025 assert(val->const_type == glsl_bool_type());
1026 uint32_t int_val =
1027 get_specialization(b, val, (opcode == SpvOpSpecConstantTrue));
1028 val->constant->values[0].u32[0] = int_val ? NIR_TRUE : NIR_FALSE;
1029 break;
1030 }
1031
1032 case SpvOpConstant: {
1033 assert(glsl_type_is_scalar(val->const_type));
1034 int bit_size = glsl_get_bit_size(val->const_type);
1035 if (bit_size == 64) {
1036 val->constant->values->u32[0] = w[3];
1037 val->constant->values->u32[1] = w[4];
1038 } else {
1039 assert(bit_size == 32);
1040 val->constant->values->u32[0] = w[3];
1041 }
1042 break;
1043 }
1044 case SpvOpSpecConstant: {
1045 assert(glsl_type_is_scalar(val->const_type));
1046 val->constant->values[0].u32[0] = get_specialization(b, val, w[3]);
1047 int bit_size = glsl_get_bit_size(val->const_type);
1048 if (bit_size == 64)
1049 val->constant->values[0].u64[0] =
1050 get_specialization64(b, val, vtn_u64_literal(&w[3]));
1051 else
1052 val->constant->values[0].u32[0] = get_specialization(b, val, w[3]);
1053 break;
1054 }
1055 case SpvOpSpecConstantComposite:
1056 case SpvOpConstantComposite: {
1057 unsigned elem_count = count - 3;
1058 nir_constant **elems = ralloc_array(b, nir_constant *, elem_count);
1059 for (unsigned i = 0; i < elem_count; i++)
1060 elems[i] = vtn_value(b, w[i + 3], vtn_value_type_constant)->constant;
1061
1062 switch (glsl_get_base_type(val->const_type)) {
1063 case GLSL_TYPE_UINT:
1064 case GLSL_TYPE_INT:
1065 case GLSL_TYPE_FLOAT:
1066 case GLSL_TYPE_BOOL:
1067 case GLSL_TYPE_DOUBLE: {
1068 int bit_size = glsl_get_bit_size(val->const_type);
1069 if (glsl_type_is_matrix(val->const_type)) {
1070 assert(glsl_get_matrix_columns(val->const_type) == elem_count);
1071 for (unsigned i = 0; i < elem_count; i++)
1072 val->constant->values[i] = elems[i]->values[0];
1073 } else {
1074 assert(glsl_type_is_vector(val->const_type));
1075 assert(glsl_get_vector_elements(val->const_type) == elem_count);
1076 for (unsigned i = 0; i < elem_count; i++) {
1077 if (bit_size == 64) {
1078 val->constant->values[0].u64[i] = elems[i]->values[0].u64[0];
1079 } else {
1080 assert(bit_size == 32);
1081 val->constant->values[0].u32[i] = elems[i]->values[0].u32[0];
1082 }
1083 }
1084 }
1085 ralloc_free(elems);
1086 break;
1087 }
1088 case GLSL_TYPE_STRUCT:
1089 case GLSL_TYPE_ARRAY:
1090 ralloc_steal(val->constant, elems);
1091 val->constant->num_elements = elem_count;
1092 val->constant->elements = elems;
1093 break;
1094
1095 default:
1096 unreachable("Unsupported type for constants");
1097 }
1098 break;
1099 }
1100
1101 case SpvOpSpecConstantOp: {
1102 SpvOp opcode = get_specialization(b, val, w[3]);
1103 switch (opcode) {
1104 case SpvOpVectorShuffle: {
1105 struct vtn_value *v0 = &b->values[w[4]];
1106 struct vtn_value *v1 = &b->values[w[5]];
1107
1108 assert(v0->value_type == vtn_value_type_constant ||
1109 v0->value_type == vtn_value_type_undef);
1110 assert(v1->value_type == vtn_value_type_constant ||
1111 v1->value_type == vtn_value_type_undef);
1112
1113 unsigned len0 = v0->value_type == vtn_value_type_constant ?
1114 glsl_get_vector_elements(v0->const_type) :
1115 glsl_get_vector_elements(v0->type->type);
1116 unsigned len1 = v1->value_type == vtn_value_type_constant ?
1117 glsl_get_vector_elements(v1->const_type) :
1118 glsl_get_vector_elements(v1->type->type);
1119
1120 assert(len0 + len1 < 16);
1121
1122 unsigned bit_size = glsl_get_bit_size(val->const_type);
1123 unsigned bit_size0 = v0->value_type == vtn_value_type_constant ?
1124 glsl_get_bit_size(v0->const_type) :
1125 glsl_get_bit_size(v0->type->type);
1126 unsigned bit_size1 = v1->value_type == vtn_value_type_constant ?
1127 glsl_get_bit_size(v1->const_type) :
1128 glsl_get_bit_size(v1->type->type);
1129
1130 assert(bit_size == bit_size0 && bit_size == bit_size1);
1131
1132 if (bit_size == 64) {
1133 uint64_t u64[8];
1134 if (v0->value_type == vtn_value_type_constant) {
1135 for (unsigned i = 0; i < len0; i++)
1136 u64[i] = v0->constant->values[0].u64[i];
1137 }
1138 if (v1->value_type == vtn_value_type_constant) {
1139 for (unsigned i = 0; i < len1; i++)
1140 u64[len0 + i] = v1->constant->values[0].u64[i];
1141 }
1142
1143 for (unsigned i = 0, j = 0; i < count - 6; i++, j++) {
1144 uint32_t comp = w[i + 6];
1145 /* If component is not used, set the value to a known constant
1146 * to detect if it is wrongly used.
1147 */
1148 if (comp == (uint32_t)-1)
1149 val->constant->values[0].u64[j] = 0xdeadbeefdeadbeef;
1150 else
1151 val->constant->values[0].u64[j] = u64[comp];
1152 }
1153 } else {
1154 uint32_t u32[8];
1155 if (v0->value_type == vtn_value_type_constant) {
1156 for (unsigned i = 0; i < len0; i++)
1157 u32[i] = v0->constant->values[0].u32[i];
1158 }
1159 if (v1->value_type == vtn_value_type_constant) {
1160 for (unsigned i = 0; i < len1; i++)
1161 u32[len0 + i] = v1->constant->values[0].u32[i];
1162 }
1163
1164 for (unsigned i = 0, j = 0; i < count - 6; i++, j++) {
1165 uint32_t comp = w[i + 6];
1166 /* If component is not used, set the value to a known constant
1167 * to detect if it is wrongly used.
1168 */
1169 if (comp == (uint32_t)-1)
1170 val->constant->values[0].u32[j] = 0xdeadbeef;
1171 else
1172 val->constant->values[0].u32[j] = u32[comp];
1173 }
1174 }
1175 break;
1176 }
1177
1178 case SpvOpCompositeExtract:
1179 case SpvOpCompositeInsert: {
1180 struct vtn_value *comp;
1181 unsigned deref_start;
1182 struct nir_constant **c;
1183 if (opcode == SpvOpCompositeExtract) {
1184 comp = vtn_value(b, w[4], vtn_value_type_constant);
1185 deref_start = 5;
1186 c = &comp->constant;
1187 } else {
1188 comp = vtn_value(b, w[5], vtn_value_type_constant);
1189 deref_start = 6;
1190 val->constant = nir_constant_clone(comp->constant,
1191 (nir_variable *)b);
1192 c = &val->constant;
1193 }
1194
1195 int elem = -1;
1196 int col = 0;
1197 const struct glsl_type *type = comp->const_type;
1198 for (unsigned i = deref_start; i < count; i++) {
1199 switch (glsl_get_base_type(type)) {
1200 case GLSL_TYPE_UINT:
1201 case GLSL_TYPE_INT:
1202 case GLSL_TYPE_FLOAT:
1203 case GLSL_TYPE_DOUBLE:
1204 case GLSL_TYPE_BOOL:
1205 /* If we hit this granularity, we're picking off an element */
1206 if (glsl_type_is_matrix(type)) {
1207 assert(col == 0 && elem == -1);
1208 col = w[i];
1209 elem = 0;
1210 type = glsl_get_column_type(type);
1211 } else {
1212 assert(elem <= 0 && glsl_type_is_vector(type));
1213 elem = w[i];
1214 type = glsl_scalar_type(glsl_get_base_type(type));
1215 }
1216 continue;
1217
1218 case GLSL_TYPE_ARRAY:
1219 c = &(*c)->elements[w[i]];
1220 type = glsl_get_array_element(type);
1221 continue;
1222
1223 case GLSL_TYPE_STRUCT:
1224 c = &(*c)->elements[w[i]];
1225 type = glsl_get_struct_field(type, w[i]);
1226 continue;
1227
1228 default:
1229 unreachable("Invalid constant type");
1230 }
1231 }
1232
1233 if (opcode == SpvOpCompositeExtract) {
1234 if (elem == -1) {
1235 val->constant = *c;
1236 } else {
1237 unsigned num_components = glsl_get_vector_elements(type);
1238 unsigned bit_size = glsl_get_bit_size(type);
1239 for (unsigned i = 0; i < num_components; i++)
1240 if (bit_size == 64) {
1241 val->constant->values[0].u64[i] = (*c)->values[col].u64[elem + i];
1242 } else {
1243 assert(bit_size == 32);
1244 val->constant->values[0].u32[i] = (*c)->values[col].u32[elem + i];
1245 }
1246 }
1247 } else {
1248 struct vtn_value *insert =
1249 vtn_value(b, w[4], vtn_value_type_constant);
1250 assert(insert->const_type == type);
1251 if (elem == -1) {
1252 *c = insert->constant;
1253 } else {
1254 unsigned num_components = glsl_get_vector_elements(type);
1255 unsigned bit_size = glsl_get_bit_size(type);
1256 for (unsigned i = 0; i < num_components; i++)
1257 if (bit_size == 64) {
1258 (*c)->values[col].u64[elem + i] = insert->constant->values[0].u64[i];
1259 } else {
1260 assert(bit_size == 32);
1261 (*c)->values[col].u32[elem + i] = insert->constant->values[0].u32[i];
1262 }
1263 }
1264 }
1265 break;
1266 }
1267
1268 default: {
1269 bool swap;
1270 nir_alu_type dst_alu_type = nir_get_nir_type_for_glsl_type(val->const_type);
1271 nir_alu_type src_alu_type = dst_alu_type;
1272 nir_op op = vtn_nir_alu_op_for_spirv_opcode(opcode, &swap, src_alu_type, dst_alu_type);
1273
1274 unsigned num_components = glsl_get_vector_elements(val->const_type);
1275 unsigned bit_size =
1276 glsl_get_bit_size(val->const_type);
1277
1278 nir_const_value src[4];
1279 assert(count <= 7);
1280 for (unsigned i = 0; i < count - 4; i++) {
1281 nir_constant *c =
1282 vtn_value(b, w[4 + i], vtn_value_type_constant)->constant;
1283
1284 unsigned j = swap ? 1 - i : i;
1285 assert(bit_size == 32);
1286 src[j] = c->values[0];
1287 }
1288
1289 val->constant->values[0] =
1290 nir_eval_const_opcode(op, num_components, bit_size, src);
1291 break;
1292 } /* default */
1293 }
1294 break;
1295 }
1296
1297 case SpvOpConstantNull:
1298 val->constant = vtn_null_constant(b, val->const_type);
1299 break;
1300
1301 case SpvOpConstantSampler:
1302 assert(!"OpConstantSampler requires Kernel Capability");
1303 break;
1304
1305 default:
1306 unreachable("Unhandled opcode");
1307 }
1308
1309 /* Now that we have the value, update the workgroup size if needed */
1310 vtn_foreach_decoration(b, val, handle_workgroup_size_decoration_cb, NULL);
1311 }
1312
1313 static void
vtn_handle_function_call(struct vtn_builder * b,SpvOp opcode,const uint32_t * w,unsigned count)1314 vtn_handle_function_call(struct vtn_builder *b, SpvOp opcode,
1315 const uint32_t *w, unsigned count)
1316 {
1317 struct nir_function *callee =
1318 vtn_value(b, w[3], vtn_value_type_function)->func->impl->function;
1319
1320 nir_call_instr *call = nir_call_instr_create(b->nb.shader, callee);
1321 for (unsigned i = 0; i < call->num_params; i++) {
1322 unsigned arg_id = w[4 + i];
1323 struct vtn_value *arg = vtn_untyped_value(b, arg_id);
1324 if (arg->value_type == vtn_value_type_access_chain) {
1325 nir_deref_var *d = vtn_access_chain_to_deref(b, arg->access_chain);
1326 call->params[i] = nir_deref_var_clone(d, call);
1327 } else {
1328 struct vtn_ssa_value *arg_ssa = vtn_ssa_value(b, arg_id);
1329
1330 /* Make a temporary to store the argument in */
1331 nir_variable *tmp =
1332 nir_local_variable_create(b->impl, arg_ssa->type, "arg_tmp");
1333 call->params[i] = nir_deref_var_create(call, tmp);
1334
1335 vtn_local_store(b, arg_ssa, call->params[i]);
1336 }
1337 }
1338
1339 nir_variable *out_tmp = NULL;
1340 if (!glsl_type_is_void(callee->return_type)) {
1341 out_tmp = nir_local_variable_create(b->impl, callee->return_type,
1342 "out_tmp");
1343 call->return_deref = nir_deref_var_create(call, out_tmp);
1344 }
1345
1346 nir_builder_instr_insert(&b->nb, &call->instr);
1347
1348 if (glsl_type_is_void(callee->return_type)) {
1349 vtn_push_value(b, w[2], vtn_value_type_undef);
1350 } else {
1351 struct vtn_value *retval = vtn_push_value(b, w[2], vtn_value_type_ssa);
1352 retval->ssa = vtn_local_load(b, call->return_deref);
1353 }
1354 }
1355
1356 struct vtn_ssa_value *
vtn_create_ssa_value(struct vtn_builder * b,const struct glsl_type * type)1357 vtn_create_ssa_value(struct vtn_builder *b, const struct glsl_type *type)
1358 {
1359 struct vtn_ssa_value *val = rzalloc(b, struct vtn_ssa_value);
1360 val->type = type;
1361
1362 if (!glsl_type_is_vector_or_scalar(type)) {
1363 unsigned elems = glsl_get_length(type);
1364 val->elems = ralloc_array(b, struct vtn_ssa_value *, elems);
1365 for (unsigned i = 0; i < elems; i++) {
1366 const struct glsl_type *child_type;
1367
1368 switch (glsl_get_base_type(type)) {
1369 case GLSL_TYPE_INT:
1370 case GLSL_TYPE_UINT:
1371 case GLSL_TYPE_BOOL:
1372 case GLSL_TYPE_FLOAT:
1373 case GLSL_TYPE_DOUBLE:
1374 child_type = glsl_get_column_type(type);
1375 break;
1376 case GLSL_TYPE_ARRAY:
1377 child_type = glsl_get_array_element(type);
1378 break;
1379 case GLSL_TYPE_STRUCT:
1380 child_type = glsl_get_struct_field(type, i);
1381 break;
1382 default:
1383 unreachable("unkown base type");
1384 }
1385
1386 val->elems[i] = vtn_create_ssa_value(b, child_type);
1387 }
1388 }
1389
1390 return val;
1391 }
1392
1393 static nir_tex_src
vtn_tex_src(struct vtn_builder * b,unsigned index,nir_tex_src_type type)1394 vtn_tex_src(struct vtn_builder *b, unsigned index, nir_tex_src_type type)
1395 {
1396 nir_tex_src src;
1397 src.src = nir_src_for_ssa(vtn_ssa_value(b, index)->def);
1398 src.src_type = type;
1399 return src;
1400 }
1401
1402 static void
vtn_handle_texture(struct vtn_builder * b,SpvOp opcode,const uint32_t * w,unsigned count)1403 vtn_handle_texture(struct vtn_builder *b, SpvOp opcode,
1404 const uint32_t *w, unsigned count)
1405 {
1406 if (opcode == SpvOpSampledImage) {
1407 struct vtn_value *val =
1408 vtn_push_value(b, w[2], vtn_value_type_sampled_image);
1409 val->sampled_image = ralloc(b, struct vtn_sampled_image);
1410 val->sampled_image->image =
1411 vtn_value(b, w[3], vtn_value_type_access_chain)->access_chain;
1412 val->sampled_image->sampler =
1413 vtn_value(b, w[4], vtn_value_type_access_chain)->access_chain;
1414 return;
1415 } else if (opcode == SpvOpImage) {
1416 struct vtn_value *val =
1417 vtn_push_value(b, w[2], vtn_value_type_access_chain);
1418 struct vtn_value *src_val = vtn_untyped_value(b, w[3]);
1419 if (src_val->value_type == vtn_value_type_sampled_image) {
1420 val->access_chain = src_val->sampled_image->image;
1421 } else {
1422 assert(src_val->value_type == vtn_value_type_access_chain);
1423 val->access_chain = src_val->access_chain;
1424 }
1425 return;
1426 }
1427
1428 struct vtn_type *ret_type = vtn_value(b, w[1], vtn_value_type_type)->type;
1429 struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_ssa);
1430
1431 struct vtn_sampled_image sampled;
1432 struct vtn_value *sampled_val = vtn_untyped_value(b, w[3]);
1433 if (sampled_val->value_type == vtn_value_type_sampled_image) {
1434 sampled = *sampled_val->sampled_image;
1435 } else {
1436 assert(sampled_val->value_type == vtn_value_type_access_chain);
1437 sampled.image = NULL;
1438 sampled.sampler = sampled_val->access_chain;
1439 }
1440
1441 const struct glsl_type *image_type;
1442 if (sampled.image) {
1443 image_type = sampled.image->var->var->interface_type;
1444 } else {
1445 image_type = sampled.sampler->var->var->interface_type;
1446 }
1447 const enum glsl_sampler_dim sampler_dim = glsl_get_sampler_dim(image_type);
1448 const bool is_array = glsl_sampler_type_is_array(image_type);
1449 const bool is_shadow = glsl_sampler_type_is_shadow(image_type);
1450
1451 /* Figure out the base texture operation */
1452 nir_texop texop;
1453 switch (opcode) {
1454 case SpvOpImageSampleImplicitLod:
1455 case SpvOpImageSampleDrefImplicitLod:
1456 case SpvOpImageSampleProjImplicitLod:
1457 case SpvOpImageSampleProjDrefImplicitLod:
1458 texop = nir_texop_tex;
1459 break;
1460
1461 case SpvOpImageSampleExplicitLod:
1462 case SpvOpImageSampleDrefExplicitLod:
1463 case SpvOpImageSampleProjExplicitLod:
1464 case SpvOpImageSampleProjDrefExplicitLod:
1465 texop = nir_texop_txl;
1466 break;
1467
1468 case SpvOpImageFetch:
1469 if (glsl_get_sampler_dim(image_type) == GLSL_SAMPLER_DIM_MS) {
1470 texop = nir_texop_txf_ms;
1471 } else {
1472 texop = nir_texop_txf;
1473 }
1474 break;
1475
1476 case SpvOpImageGather:
1477 case SpvOpImageDrefGather:
1478 texop = nir_texop_tg4;
1479 break;
1480
1481 case SpvOpImageQuerySizeLod:
1482 case SpvOpImageQuerySize:
1483 texop = nir_texop_txs;
1484 break;
1485
1486 case SpvOpImageQueryLod:
1487 texop = nir_texop_lod;
1488 break;
1489
1490 case SpvOpImageQueryLevels:
1491 texop = nir_texop_query_levels;
1492 break;
1493
1494 case SpvOpImageQuerySamples:
1495 texop = nir_texop_texture_samples;
1496 break;
1497
1498 default:
1499 unreachable("Unhandled opcode");
1500 }
1501
1502 nir_tex_src srcs[8]; /* 8 should be enough */
1503 nir_tex_src *p = srcs;
1504
1505 unsigned idx = 4;
1506
1507 struct nir_ssa_def *coord;
1508 unsigned coord_components;
1509 switch (opcode) {
1510 case SpvOpImageSampleImplicitLod:
1511 case SpvOpImageSampleExplicitLod:
1512 case SpvOpImageSampleDrefImplicitLod:
1513 case SpvOpImageSampleDrefExplicitLod:
1514 case SpvOpImageSampleProjImplicitLod:
1515 case SpvOpImageSampleProjExplicitLod:
1516 case SpvOpImageSampleProjDrefImplicitLod:
1517 case SpvOpImageSampleProjDrefExplicitLod:
1518 case SpvOpImageFetch:
1519 case SpvOpImageGather:
1520 case SpvOpImageDrefGather:
1521 case SpvOpImageQueryLod: {
1522 /* All these types have the coordinate as their first real argument */
1523 switch (sampler_dim) {
1524 case GLSL_SAMPLER_DIM_1D:
1525 case GLSL_SAMPLER_DIM_BUF:
1526 coord_components = 1;
1527 break;
1528 case GLSL_SAMPLER_DIM_2D:
1529 case GLSL_SAMPLER_DIM_RECT:
1530 case GLSL_SAMPLER_DIM_MS:
1531 coord_components = 2;
1532 break;
1533 case GLSL_SAMPLER_DIM_3D:
1534 case GLSL_SAMPLER_DIM_CUBE:
1535 coord_components = 3;
1536 break;
1537 default:
1538 unreachable("Invalid sampler type");
1539 }
1540
1541 if (is_array && texop != nir_texop_lod)
1542 coord_components++;
1543
1544 coord = vtn_ssa_value(b, w[idx++])->def;
1545 p->src = nir_src_for_ssa(coord);
1546 p->src_type = nir_tex_src_coord;
1547 p++;
1548 break;
1549 }
1550
1551 default:
1552 coord = NULL;
1553 coord_components = 0;
1554 break;
1555 }
1556
1557 switch (opcode) {
1558 case SpvOpImageSampleProjImplicitLod:
1559 case SpvOpImageSampleProjExplicitLod:
1560 case SpvOpImageSampleProjDrefImplicitLod:
1561 case SpvOpImageSampleProjDrefExplicitLod:
1562 /* These have the projector as the last coordinate component */
1563 p->src = nir_src_for_ssa(nir_channel(&b->nb, coord, coord_components));
1564 p->src_type = nir_tex_src_projector;
1565 p++;
1566 break;
1567
1568 default:
1569 break;
1570 }
1571
1572 unsigned gather_component = 0;
1573 switch (opcode) {
1574 case SpvOpImageSampleDrefImplicitLod:
1575 case SpvOpImageSampleDrefExplicitLod:
1576 case SpvOpImageSampleProjDrefImplicitLod:
1577 case SpvOpImageSampleProjDrefExplicitLod:
1578 case SpvOpImageDrefGather:
1579 /* These all have an explicit depth value as their next source */
1580 (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_comparator);
1581 break;
1582
1583 case SpvOpImageGather:
1584 /* This has a component as its next source */
1585 gather_component =
1586 vtn_value(b, w[idx++], vtn_value_type_constant)->constant->values[0].u32[0];
1587 break;
1588
1589 default:
1590 break;
1591 }
1592
1593 /* For OpImageQuerySizeLod, we always have an LOD */
1594 if (opcode == SpvOpImageQuerySizeLod)
1595 (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_lod);
1596
1597 /* Now we need to handle some number of optional arguments */
1598 const struct vtn_ssa_value *gather_offsets = NULL;
1599 if (idx < count) {
1600 uint32_t operands = w[idx++];
1601
1602 if (operands & SpvImageOperandsBiasMask) {
1603 assert(texop == nir_texop_tex);
1604 texop = nir_texop_txb;
1605 (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_bias);
1606 }
1607
1608 if (operands & SpvImageOperandsLodMask) {
1609 assert(texop == nir_texop_txl || texop == nir_texop_txf ||
1610 texop == nir_texop_txs);
1611 (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_lod);
1612 }
1613
1614 if (operands & SpvImageOperandsGradMask) {
1615 assert(texop == nir_texop_txl);
1616 texop = nir_texop_txd;
1617 (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_ddx);
1618 (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_ddy);
1619 }
1620
1621 if (operands & SpvImageOperandsOffsetMask ||
1622 operands & SpvImageOperandsConstOffsetMask)
1623 (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_offset);
1624
1625 if (operands & SpvImageOperandsConstOffsetsMask) {
1626 gather_offsets = vtn_ssa_value(b, w[idx++]);
1627 (*p++) = (nir_tex_src){};
1628 }
1629
1630 if (operands & SpvImageOperandsSampleMask) {
1631 assert(texop == nir_texop_txf_ms);
1632 texop = nir_texop_txf_ms;
1633 (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_ms_index);
1634 }
1635 }
1636 /* We should have now consumed exactly all of the arguments */
1637 assert(idx == count);
1638
1639 nir_tex_instr *instr = nir_tex_instr_create(b->shader, p - srcs);
1640 instr->op = texop;
1641
1642 memcpy(instr->src, srcs, instr->num_srcs * sizeof(*instr->src));
1643
1644 instr->coord_components = coord_components;
1645 instr->sampler_dim = sampler_dim;
1646 instr->is_array = is_array;
1647 instr->is_shadow = is_shadow;
1648 instr->is_new_style_shadow =
1649 is_shadow && glsl_get_components(ret_type->type) == 1;
1650 instr->component = gather_component;
1651
1652 switch (glsl_get_sampler_result_type(image_type)) {
1653 case GLSL_TYPE_FLOAT: instr->dest_type = nir_type_float; break;
1654 case GLSL_TYPE_INT: instr->dest_type = nir_type_int; break;
1655 case GLSL_TYPE_UINT: instr->dest_type = nir_type_uint; break;
1656 case GLSL_TYPE_BOOL: instr->dest_type = nir_type_bool; break;
1657 default:
1658 unreachable("Invalid base type for sampler result");
1659 }
1660
1661 nir_deref_var *sampler = vtn_access_chain_to_deref(b, sampled.sampler);
1662 nir_deref_var *texture;
1663 if (sampled.image) {
1664 nir_deref_var *image = vtn_access_chain_to_deref(b, sampled.image);
1665 texture = image;
1666 } else {
1667 texture = sampler;
1668 }
1669
1670 instr->texture = nir_deref_var_clone(texture, instr);
1671
1672 switch (instr->op) {
1673 case nir_texop_tex:
1674 case nir_texop_txb:
1675 case nir_texop_txl:
1676 case nir_texop_txd:
1677 /* These operations require a sampler */
1678 instr->sampler = nir_deref_var_clone(sampler, instr);
1679 break;
1680 case nir_texop_txf:
1681 case nir_texop_txf_ms:
1682 case nir_texop_txs:
1683 case nir_texop_lod:
1684 case nir_texop_tg4:
1685 case nir_texop_query_levels:
1686 case nir_texop_texture_samples:
1687 case nir_texop_samples_identical:
1688 /* These don't */
1689 instr->sampler = NULL;
1690 break;
1691 case nir_texop_txf_ms_mcs:
1692 unreachable("unexpected nir_texop_txf_ms_mcs");
1693 }
1694
1695 nir_ssa_dest_init(&instr->instr, &instr->dest,
1696 nir_tex_instr_dest_size(instr), 32, NULL);
1697
1698 assert(glsl_get_vector_elements(ret_type->type) ==
1699 nir_tex_instr_dest_size(instr));
1700
1701 nir_ssa_def *def;
1702 nir_instr *instruction;
1703 if (gather_offsets) {
1704 assert(glsl_get_base_type(gather_offsets->type) == GLSL_TYPE_ARRAY);
1705 assert(glsl_get_length(gather_offsets->type) == 4);
1706 nir_tex_instr *instrs[4] = {instr, NULL, NULL, NULL};
1707
1708 /* Copy the current instruction 4x */
1709 for (uint32_t i = 1; i < 4; i++) {
1710 instrs[i] = nir_tex_instr_create(b->shader, instr->num_srcs);
1711 instrs[i]->op = instr->op;
1712 instrs[i]->coord_components = instr->coord_components;
1713 instrs[i]->sampler_dim = instr->sampler_dim;
1714 instrs[i]->is_array = instr->is_array;
1715 instrs[i]->is_shadow = instr->is_shadow;
1716 instrs[i]->is_new_style_shadow = instr->is_new_style_shadow;
1717 instrs[i]->component = instr->component;
1718 instrs[i]->dest_type = instr->dest_type;
1719 instrs[i]->texture = nir_deref_var_clone(texture, instrs[i]);
1720 instrs[i]->sampler = NULL;
1721
1722 memcpy(instrs[i]->src, srcs, instr->num_srcs * sizeof(*instr->src));
1723
1724 nir_ssa_dest_init(&instrs[i]->instr, &instrs[i]->dest,
1725 nir_tex_instr_dest_size(instr), 32, NULL);
1726 }
1727
1728 /* Fill in the last argument with the offset from the passed in offsets
1729 * and insert the instruction into the stream.
1730 */
1731 for (uint32_t i = 0; i < 4; i++) {
1732 nir_tex_src src;
1733 src.src = nir_src_for_ssa(gather_offsets->elems[i]->def);
1734 src.src_type = nir_tex_src_offset;
1735 instrs[i]->src[instrs[i]->num_srcs - 1] = src;
1736 nir_builder_instr_insert(&b->nb, &instrs[i]->instr);
1737 }
1738
1739 /* Combine the results of the 4 instructions by taking their .w
1740 * components
1741 */
1742 nir_alu_instr *vec4 = nir_alu_instr_create(b->shader, nir_op_vec4);
1743 nir_ssa_dest_init(&vec4->instr, &vec4->dest.dest, 4, 32, NULL);
1744 vec4->dest.write_mask = 0xf;
1745 for (uint32_t i = 0; i < 4; i++) {
1746 vec4->src[i].src = nir_src_for_ssa(&instrs[i]->dest.ssa);
1747 vec4->src[i].swizzle[0] = 3;
1748 }
1749 def = &vec4->dest.dest.ssa;
1750 instruction = &vec4->instr;
1751 } else {
1752 def = &instr->dest.ssa;
1753 instruction = &instr->instr;
1754 }
1755
1756 val->ssa = vtn_create_ssa_value(b, ret_type->type);
1757 val->ssa->def = def;
1758
1759 nir_builder_instr_insert(&b->nb, instruction);
1760 }
1761
1762 static void
fill_common_atomic_sources(struct vtn_builder * b,SpvOp opcode,const uint32_t * w,nir_src * src)1763 fill_common_atomic_sources(struct vtn_builder *b, SpvOp opcode,
1764 const uint32_t *w, nir_src *src)
1765 {
1766 switch (opcode) {
1767 case SpvOpAtomicIIncrement:
1768 src[0] = nir_src_for_ssa(nir_imm_int(&b->nb, 1));
1769 break;
1770
1771 case SpvOpAtomicIDecrement:
1772 src[0] = nir_src_for_ssa(nir_imm_int(&b->nb, -1));
1773 break;
1774
1775 case SpvOpAtomicISub:
1776 src[0] =
1777 nir_src_for_ssa(nir_ineg(&b->nb, vtn_ssa_value(b, w[6])->def));
1778 break;
1779
1780 case SpvOpAtomicCompareExchange:
1781 src[0] = nir_src_for_ssa(vtn_ssa_value(b, w[8])->def);
1782 src[1] = nir_src_for_ssa(vtn_ssa_value(b, w[7])->def);
1783 break;
1784
1785 case SpvOpAtomicExchange:
1786 case SpvOpAtomicIAdd:
1787 case SpvOpAtomicSMin:
1788 case SpvOpAtomicUMin:
1789 case SpvOpAtomicSMax:
1790 case SpvOpAtomicUMax:
1791 case SpvOpAtomicAnd:
1792 case SpvOpAtomicOr:
1793 case SpvOpAtomicXor:
1794 src[0] = nir_src_for_ssa(vtn_ssa_value(b, w[6])->def);
1795 break;
1796
1797 default:
1798 unreachable("Invalid SPIR-V atomic");
1799 }
1800 }
1801
1802 static nir_ssa_def *
get_image_coord(struct vtn_builder * b,uint32_t value)1803 get_image_coord(struct vtn_builder *b, uint32_t value)
1804 {
1805 struct vtn_ssa_value *coord = vtn_ssa_value(b, value);
1806
1807 /* The image_load_store intrinsics assume a 4-dim coordinate */
1808 unsigned dim = glsl_get_vector_elements(coord->type);
1809 unsigned swizzle[4];
1810 for (unsigned i = 0; i < 4; i++)
1811 swizzle[i] = MIN2(i, dim - 1);
1812
1813 return nir_swizzle(&b->nb, coord->def, swizzle, 4, false);
1814 }
1815
1816 static void
vtn_handle_image(struct vtn_builder * b,SpvOp opcode,const uint32_t * w,unsigned count)1817 vtn_handle_image(struct vtn_builder *b, SpvOp opcode,
1818 const uint32_t *w, unsigned count)
1819 {
1820 /* Just get this one out of the way */
1821 if (opcode == SpvOpImageTexelPointer) {
1822 struct vtn_value *val =
1823 vtn_push_value(b, w[2], vtn_value_type_image_pointer);
1824 val->image = ralloc(b, struct vtn_image_pointer);
1825
1826 val->image->image =
1827 vtn_value(b, w[3], vtn_value_type_access_chain)->access_chain;
1828 val->image->coord = get_image_coord(b, w[4]);
1829 val->image->sample = vtn_ssa_value(b, w[5])->def;
1830 return;
1831 }
1832
1833 struct vtn_image_pointer image;
1834
1835 switch (opcode) {
1836 case SpvOpAtomicExchange:
1837 case SpvOpAtomicCompareExchange:
1838 case SpvOpAtomicCompareExchangeWeak:
1839 case SpvOpAtomicIIncrement:
1840 case SpvOpAtomicIDecrement:
1841 case SpvOpAtomicIAdd:
1842 case SpvOpAtomicISub:
1843 case SpvOpAtomicLoad:
1844 case SpvOpAtomicSMin:
1845 case SpvOpAtomicUMin:
1846 case SpvOpAtomicSMax:
1847 case SpvOpAtomicUMax:
1848 case SpvOpAtomicAnd:
1849 case SpvOpAtomicOr:
1850 case SpvOpAtomicXor:
1851 image = *vtn_value(b, w[3], vtn_value_type_image_pointer)->image;
1852 break;
1853
1854 case SpvOpAtomicStore:
1855 image = *vtn_value(b, w[1], vtn_value_type_image_pointer)->image;
1856 break;
1857
1858 case SpvOpImageQuerySize:
1859 image.image =
1860 vtn_value(b, w[3], vtn_value_type_access_chain)->access_chain;
1861 image.coord = NULL;
1862 image.sample = NULL;
1863 break;
1864
1865 case SpvOpImageRead:
1866 image.image =
1867 vtn_value(b, w[3], vtn_value_type_access_chain)->access_chain;
1868 image.coord = get_image_coord(b, w[4]);
1869
1870 if (count > 5 && (w[5] & SpvImageOperandsSampleMask)) {
1871 assert(w[5] == SpvImageOperandsSampleMask);
1872 image.sample = vtn_ssa_value(b, w[6])->def;
1873 } else {
1874 image.sample = nir_ssa_undef(&b->nb, 1, 32);
1875 }
1876 break;
1877
1878 case SpvOpImageWrite:
1879 image.image =
1880 vtn_value(b, w[1], vtn_value_type_access_chain)->access_chain;
1881 image.coord = get_image_coord(b, w[2]);
1882
1883 /* texel = w[3] */
1884
1885 if (count > 4 && (w[4] & SpvImageOperandsSampleMask)) {
1886 assert(w[4] == SpvImageOperandsSampleMask);
1887 image.sample = vtn_ssa_value(b, w[5])->def;
1888 } else {
1889 image.sample = nir_ssa_undef(&b->nb, 1, 32);
1890 }
1891 break;
1892
1893 default:
1894 unreachable("Invalid image opcode");
1895 }
1896
1897 nir_intrinsic_op op;
1898 switch (opcode) {
1899 #define OP(S, N) case SpvOp##S: op = nir_intrinsic_image_##N; break;
1900 OP(ImageQuerySize, size)
1901 OP(ImageRead, load)
1902 OP(ImageWrite, store)
1903 OP(AtomicLoad, load)
1904 OP(AtomicStore, store)
1905 OP(AtomicExchange, atomic_exchange)
1906 OP(AtomicCompareExchange, atomic_comp_swap)
1907 OP(AtomicIIncrement, atomic_add)
1908 OP(AtomicIDecrement, atomic_add)
1909 OP(AtomicIAdd, atomic_add)
1910 OP(AtomicISub, atomic_add)
1911 OP(AtomicSMin, atomic_min)
1912 OP(AtomicUMin, atomic_min)
1913 OP(AtomicSMax, atomic_max)
1914 OP(AtomicUMax, atomic_max)
1915 OP(AtomicAnd, atomic_and)
1916 OP(AtomicOr, atomic_or)
1917 OP(AtomicXor, atomic_xor)
1918 #undef OP
1919 default:
1920 unreachable("Invalid image opcode");
1921 }
1922
1923 nir_intrinsic_instr *intrin = nir_intrinsic_instr_create(b->shader, op);
1924
1925 nir_deref_var *image_deref = vtn_access_chain_to_deref(b, image.image);
1926 intrin->variables[0] = nir_deref_var_clone(image_deref, intrin);
1927
1928 /* ImageQuerySize doesn't take any extra parameters */
1929 if (opcode != SpvOpImageQuerySize) {
1930 /* The image coordinate is always 4 components but we may not have that
1931 * many. Swizzle to compensate.
1932 */
1933 unsigned swiz[4];
1934 for (unsigned i = 0; i < 4; i++)
1935 swiz[i] = i < image.coord->num_components ? i : 0;
1936 intrin->src[0] = nir_src_for_ssa(nir_swizzle(&b->nb, image.coord,
1937 swiz, 4, false));
1938 intrin->src[1] = nir_src_for_ssa(image.sample);
1939 }
1940
1941 switch (opcode) {
1942 case SpvOpAtomicLoad:
1943 case SpvOpImageQuerySize:
1944 case SpvOpImageRead:
1945 break;
1946 case SpvOpAtomicStore:
1947 intrin->src[2] = nir_src_for_ssa(vtn_ssa_value(b, w[4])->def);
1948 break;
1949 case SpvOpImageWrite:
1950 intrin->src[2] = nir_src_for_ssa(vtn_ssa_value(b, w[3])->def);
1951 break;
1952
1953 case SpvOpAtomicIIncrement:
1954 case SpvOpAtomicIDecrement:
1955 case SpvOpAtomicExchange:
1956 case SpvOpAtomicIAdd:
1957 case SpvOpAtomicSMin:
1958 case SpvOpAtomicUMin:
1959 case SpvOpAtomicSMax:
1960 case SpvOpAtomicUMax:
1961 case SpvOpAtomicAnd:
1962 case SpvOpAtomicOr:
1963 case SpvOpAtomicXor:
1964 fill_common_atomic_sources(b, opcode, w, &intrin->src[2]);
1965 break;
1966
1967 default:
1968 unreachable("Invalid image opcode");
1969 }
1970
1971 if (opcode != SpvOpImageWrite) {
1972 struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_ssa);
1973 struct vtn_type *type = vtn_value(b, w[1], vtn_value_type_type)->type;
1974 nir_ssa_dest_init(&intrin->instr, &intrin->dest, 4, 32, NULL);
1975
1976 nir_builder_instr_insert(&b->nb, &intrin->instr);
1977
1978 /* The image intrinsics always return 4 channels but we may not want
1979 * that many. Emit a mov to trim it down.
1980 */
1981 unsigned swiz[4] = {0, 1, 2, 3};
1982 val->ssa = vtn_create_ssa_value(b, type->type);
1983 val->ssa->def = nir_swizzle(&b->nb, &intrin->dest.ssa, swiz,
1984 glsl_get_vector_elements(type->type), false);
1985 } else {
1986 nir_builder_instr_insert(&b->nb, &intrin->instr);
1987 }
1988 }
1989
1990 static nir_intrinsic_op
get_ssbo_nir_atomic_op(SpvOp opcode)1991 get_ssbo_nir_atomic_op(SpvOp opcode)
1992 {
1993 switch (opcode) {
1994 case SpvOpAtomicLoad: return nir_intrinsic_load_ssbo;
1995 case SpvOpAtomicStore: return nir_intrinsic_store_ssbo;
1996 #define OP(S, N) case SpvOp##S: return nir_intrinsic_ssbo_##N;
1997 OP(AtomicExchange, atomic_exchange)
1998 OP(AtomicCompareExchange, atomic_comp_swap)
1999 OP(AtomicIIncrement, atomic_add)
2000 OP(AtomicIDecrement, atomic_add)
2001 OP(AtomicIAdd, atomic_add)
2002 OP(AtomicISub, atomic_add)
2003 OP(AtomicSMin, atomic_imin)
2004 OP(AtomicUMin, atomic_umin)
2005 OP(AtomicSMax, atomic_imax)
2006 OP(AtomicUMax, atomic_umax)
2007 OP(AtomicAnd, atomic_and)
2008 OP(AtomicOr, atomic_or)
2009 OP(AtomicXor, atomic_xor)
2010 #undef OP
2011 default:
2012 unreachable("Invalid SSBO atomic");
2013 }
2014 }
2015
2016 static nir_intrinsic_op
get_shared_nir_atomic_op(SpvOp opcode)2017 get_shared_nir_atomic_op(SpvOp opcode)
2018 {
2019 switch (opcode) {
2020 case SpvOpAtomicLoad: return nir_intrinsic_load_var;
2021 case SpvOpAtomicStore: return nir_intrinsic_store_var;
2022 #define OP(S, N) case SpvOp##S: return nir_intrinsic_var_##N;
2023 OP(AtomicExchange, atomic_exchange)
2024 OP(AtomicCompareExchange, atomic_comp_swap)
2025 OP(AtomicIIncrement, atomic_add)
2026 OP(AtomicIDecrement, atomic_add)
2027 OP(AtomicIAdd, atomic_add)
2028 OP(AtomicISub, atomic_add)
2029 OP(AtomicSMin, atomic_imin)
2030 OP(AtomicUMin, atomic_umin)
2031 OP(AtomicSMax, atomic_imax)
2032 OP(AtomicUMax, atomic_umax)
2033 OP(AtomicAnd, atomic_and)
2034 OP(AtomicOr, atomic_or)
2035 OP(AtomicXor, atomic_xor)
2036 #undef OP
2037 default:
2038 unreachable("Invalid shared atomic");
2039 }
2040 }
2041
2042 static void
vtn_handle_ssbo_or_shared_atomic(struct vtn_builder * b,SpvOp opcode,const uint32_t * w,unsigned count)2043 vtn_handle_ssbo_or_shared_atomic(struct vtn_builder *b, SpvOp opcode,
2044 const uint32_t *w, unsigned count)
2045 {
2046 struct vtn_access_chain *chain;
2047 nir_intrinsic_instr *atomic;
2048
2049 switch (opcode) {
2050 case SpvOpAtomicLoad:
2051 case SpvOpAtomicExchange:
2052 case SpvOpAtomicCompareExchange:
2053 case SpvOpAtomicCompareExchangeWeak:
2054 case SpvOpAtomicIIncrement:
2055 case SpvOpAtomicIDecrement:
2056 case SpvOpAtomicIAdd:
2057 case SpvOpAtomicISub:
2058 case SpvOpAtomicSMin:
2059 case SpvOpAtomicUMin:
2060 case SpvOpAtomicSMax:
2061 case SpvOpAtomicUMax:
2062 case SpvOpAtomicAnd:
2063 case SpvOpAtomicOr:
2064 case SpvOpAtomicXor:
2065 chain =
2066 vtn_value(b, w[3], vtn_value_type_access_chain)->access_chain;
2067 break;
2068
2069 case SpvOpAtomicStore:
2070 chain =
2071 vtn_value(b, w[1], vtn_value_type_access_chain)->access_chain;
2072 break;
2073
2074 default:
2075 unreachable("Invalid SPIR-V atomic");
2076 }
2077
2078 /*
2079 SpvScope scope = w[4];
2080 SpvMemorySemanticsMask semantics = w[5];
2081 */
2082
2083 if (chain->var->mode == vtn_variable_mode_workgroup) {
2084 struct vtn_type *type = chain->var->type;
2085 nir_deref_var *deref = vtn_access_chain_to_deref(b, chain);
2086 nir_intrinsic_op op = get_shared_nir_atomic_op(opcode);
2087 atomic = nir_intrinsic_instr_create(b->nb.shader, op);
2088 atomic->variables[0] = nir_deref_var_clone(deref, atomic);
2089
2090 switch (opcode) {
2091 case SpvOpAtomicLoad:
2092 atomic->num_components = glsl_get_vector_elements(type->type);
2093 break;
2094
2095 case SpvOpAtomicStore:
2096 atomic->num_components = glsl_get_vector_elements(type->type);
2097 nir_intrinsic_set_write_mask(atomic, (1 << atomic->num_components) - 1);
2098 atomic->src[0] = nir_src_for_ssa(vtn_ssa_value(b, w[4])->def);
2099 break;
2100
2101 case SpvOpAtomicExchange:
2102 case SpvOpAtomicCompareExchange:
2103 case SpvOpAtomicCompareExchangeWeak:
2104 case SpvOpAtomicIIncrement:
2105 case SpvOpAtomicIDecrement:
2106 case SpvOpAtomicIAdd:
2107 case SpvOpAtomicISub:
2108 case SpvOpAtomicSMin:
2109 case SpvOpAtomicUMin:
2110 case SpvOpAtomicSMax:
2111 case SpvOpAtomicUMax:
2112 case SpvOpAtomicAnd:
2113 case SpvOpAtomicOr:
2114 case SpvOpAtomicXor:
2115 fill_common_atomic_sources(b, opcode, w, &atomic->src[0]);
2116 break;
2117
2118 default:
2119 unreachable("Invalid SPIR-V atomic");
2120
2121 }
2122 } else {
2123 assert(chain->var->mode == vtn_variable_mode_ssbo);
2124 struct vtn_type *type;
2125 nir_ssa_def *offset, *index;
2126 offset = vtn_access_chain_to_offset(b, chain, &index, &type, NULL, false);
2127
2128 nir_intrinsic_op op = get_ssbo_nir_atomic_op(opcode);
2129
2130 atomic = nir_intrinsic_instr_create(b->nb.shader, op);
2131
2132 switch (opcode) {
2133 case SpvOpAtomicLoad:
2134 atomic->num_components = glsl_get_vector_elements(type->type);
2135 atomic->src[0] = nir_src_for_ssa(index);
2136 atomic->src[1] = nir_src_for_ssa(offset);
2137 break;
2138
2139 case SpvOpAtomicStore:
2140 atomic->num_components = glsl_get_vector_elements(type->type);
2141 nir_intrinsic_set_write_mask(atomic, (1 << atomic->num_components) - 1);
2142 atomic->src[0] = nir_src_for_ssa(vtn_ssa_value(b, w[4])->def);
2143 atomic->src[1] = nir_src_for_ssa(index);
2144 atomic->src[2] = nir_src_for_ssa(offset);
2145 break;
2146
2147 case SpvOpAtomicExchange:
2148 case SpvOpAtomicCompareExchange:
2149 case SpvOpAtomicCompareExchangeWeak:
2150 case SpvOpAtomicIIncrement:
2151 case SpvOpAtomicIDecrement:
2152 case SpvOpAtomicIAdd:
2153 case SpvOpAtomicISub:
2154 case SpvOpAtomicSMin:
2155 case SpvOpAtomicUMin:
2156 case SpvOpAtomicSMax:
2157 case SpvOpAtomicUMax:
2158 case SpvOpAtomicAnd:
2159 case SpvOpAtomicOr:
2160 case SpvOpAtomicXor:
2161 atomic->src[0] = nir_src_for_ssa(index);
2162 atomic->src[1] = nir_src_for_ssa(offset);
2163 fill_common_atomic_sources(b, opcode, w, &atomic->src[2]);
2164 break;
2165
2166 default:
2167 unreachable("Invalid SPIR-V atomic");
2168 }
2169 }
2170
2171 if (opcode != SpvOpAtomicStore) {
2172 struct vtn_type *type = vtn_value(b, w[1], vtn_value_type_type)->type;
2173
2174 nir_ssa_dest_init(&atomic->instr, &atomic->dest,
2175 glsl_get_vector_elements(type->type),
2176 glsl_get_bit_size(type->type), NULL);
2177
2178 struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_ssa);
2179 val->ssa = rzalloc(b, struct vtn_ssa_value);
2180 val->ssa->def = &atomic->dest.ssa;
2181 val->ssa->type = type->type;
2182 }
2183
2184 nir_builder_instr_insert(&b->nb, &atomic->instr);
2185 }
2186
2187 static nir_alu_instr *
create_vec(nir_shader * shader,unsigned num_components,unsigned bit_size)2188 create_vec(nir_shader *shader, unsigned num_components, unsigned bit_size)
2189 {
2190 nir_op op;
2191 switch (num_components) {
2192 case 1: op = nir_op_fmov; break;
2193 case 2: op = nir_op_vec2; break;
2194 case 3: op = nir_op_vec3; break;
2195 case 4: op = nir_op_vec4; break;
2196 default: unreachable("bad vector size");
2197 }
2198
2199 nir_alu_instr *vec = nir_alu_instr_create(shader, op);
2200 nir_ssa_dest_init(&vec->instr, &vec->dest.dest, num_components,
2201 bit_size, NULL);
2202 vec->dest.write_mask = (1 << num_components) - 1;
2203
2204 return vec;
2205 }
2206
2207 struct vtn_ssa_value *
vtn_ssa_transpose(struct vtn_builder * b,struct vtn_ssa_value * src)2208 vtn_ssa_transpose(struct vtn_builder *b, struct vtn_ssa_value *src)
2209 {
2210 if (src->transposed)
2211 return src->transposed;
2212
2213 struct vtn_ssa_value *dest =
2214 vtn_create_ssa_value(b, glsl_transposed_type(src->type));
2215
2216 for (unsigned i = 0; i < glsl_get_matrix_columns(dest->type); i++) {
2217 nir_alu_instr *vec = create_vec(b->shader,
2218 glsl_get_matrix_columns(src->type),
2219 glsl_get_bit_size(src->type));
2220 if (glsl_type_is_vector_or_scalar(src->type)) {
2221 vec->src[0].src = nir_src_for_ssa(src->def);
2222 vec->src[0].swizzle[0] = i;
2223 } else {
2224 for (unsigned j = 0; j < glsl_get_matrix_columns(src->type); j++) {
2225 vec->src[j].src = nir_src_for_ssa(src->elems[j]->def);
2226 vec->src[j].swizzle[0] = i;
2227 }
2228 }
2229 nir_builder_instr_insert(&b->nb, &vec->instr);
2230 dest->elems[i]->def = &vec->dest.dest.ssa;
2231 }
2232
2233 dest->transposed = src;
2234
2235 return dest;
2236 }
2237
2238 nir_ssa_def *
vtn_vector_extract(struct vtn_builder * b,nir_ssa_def * src,unsigned index)2239 vtn_vector_extract(struct vtn_builder *b, nir_ssa_def *src, unsigned index)
2240 {
2241 unsigned swiz[4] = { index };
2242 return nir_swizzle(&b->nb, src, swiz, 1, true);
2243 }
2244
2245 nir_ssa_def *
vtn_vector_insert(struct vtn_builder * b,nir_ssa_def * src,nir_ssa_def * insert,unsigned index)2246 vtn_vector_insert(struct vtn_builder *b, nir_ssa_def *src, nir_ssa_def *insert,
2247 unsigned index)
2248 {
2249 nir_alu_instr *vec = create_vec(b->shader, src->num_components,
2250 src->bit_size);
2251
2252 for (unsigned i = 0; i < src->num_components; i++) {
2253 if (i == index) {
2254 vec->src[i].src = nir_src_for_ssa(insert);
2255 } else {
2256 vec->src[i].src = nir_src_for_ssa(src);
2257 vec->src[i].swizzle[0] = i;
2258 }
2259 }
2260
2261 nir_builder_instr_insert(&b->nb, &vec->instr);
2262
2263 return &vec->dest.dest.ssa;
2264 }
2265
2266 nir_ssa_def *
vtn_vector_extract_dynamic(struct vtn_builder * b,nir_ssa_def * src,nir_ssa_def * index)2267 vtn_vector_extract_dynamic(struct vtn_builder *b, nir_ssa_def *src,
2268 nir_ssa_def *index)
2269 {
2270 nir_ssa_def *dest = vtn_vector_extract(b, src, 0);
2271 for (unsigned i = 1; i < src->num_components; i++)
2272 dest = nir_bcsel(&b->nb, nir_ieq(&b->nb, index, nir_imm_int(&b->nb, i)),
2273 vtn_vector_extract(b, src, i), dest);
2274
2275 return dest;
2276 }
2277
2278 nir_ssa_def *
vtn_vector_insert_dynamic(struct vtn_builder * b,nir_ssa_def * src,nir_ssa_def * insert,nir_ssa_def * index)2279 vtn_vector_insert_dynamic(struct vtn_builder *b, nir_ssa_def *src,
2280 nir_ssa_def *insert, nir_ssa_def *index)
2281 {
2282 nir_ssa_def *dest = vtn_vector_insert(b, src, insert, 0);
2283 for (unsigned i = 1; i < src->num_components; i++)
2284 dest = nir_bcsel(&b->nb, nir_ieq(&b->nb, index, nir_imm_int(&b->nb, i)),
2285 vtn_vector_insert(b, src, insert, i), dest);
2286
2287 return dest;
2288 }
2289
2290 static nir_ssa_def *
vtn_vector_shuffle(struct vtn_builder * b,unsigned num_components,nir_ssa_def * src0,nir_ssa_def * src1,const uint32_t * indices)2291 vtn_vector_shuffle(struct vtn_builder *b, unsigned num_components,
2292 nir_ssa_def *src0, nir_ssa_def *src1,
2293 const uint32_t *indices)
2294 {
2295 nir_alu_instr *vec = create_vec(b->shader, num_components, src0->bit_size);
2296
2297 for (unsigned i = 0; i < num_components; i++) {
2298 uint32_t index = indices[i];
2299 if (index == 0xffffffff) {
2300 vec->src[i].src =
2301 nir_src_for_ssa(nir_ssa_undef(&b->nb, 1, src0->bit_size));
2302 } else if (index < src0->num_components) {
2303 vec->src[i].src = nir_src_for_ssa(src0);
2304 vec->src[i].swizzle[0] = index;
2305 } else {
2306 vec->src[i].src = nir_src_for_ssa(src1);
2307 vec->src[i].swizzle[0] = index - src0->num_components;
2308 }
2309 }
2310
2311 nir_builder_instr_insert(&b->nb, &vec->instr);
2312
2313 return &vec->dest.dest.ssa;
2314 }
2315
2316 /*
2317 * Concatentates a number of vectors/scalars together to produce a vector
2318 */
2319 static nir_ssa_def *
vtn_vector_construct(struct vtn_builder * b,unsigned num_components,unsigned num_srcs,nir_ssa_def ** srcs)2320 vtn_vector_construct(struct vtn_builder *b, unsigned num_components,
2321 unsigned num_srcs, nir_ssa_def **srcs)
2322 {
2323 nir_alu_instr *vec = create_vec(b->shader, num_components,
2324 srcs[0]->bit_size);
2325
2326 unsigned dest_idx = 0;
2327 for (unsigned i = 0; i < num_srcs; i++) {
2328 nir_ssa_def *src = srcs[i];
2329 for (unsigned j = 0; j < src->num_components; j++) {
2330 vec->src[dest_idx].src = nir_src_for_ssa(src);
2331 vec->src[dest_idx].swizzle[0] = j;
2332 dest_idx++;
2333 }
2334 }
2335
2336 nir_builder_instr_insert(&b->nb, &vec->instr);
2337
2338 return &vec->dest.dest.ssa;
2339 }
2340
2341 static struct vtn_ssa_value *
vtn_composite_copy(void * mem_ctx,struct vtn_ssa_value * src)2342 vtn_composite_copy(void *mem_ctx, struct vtn_ssa_value *src)
2343 {
2344 struct vtn_ssa_value *dest = rzalloc(mem_ctx, struct vtn_ssa_value);
2345 dest->type = src->type;
2346
2347 if (glsl_type_is_vector_or_scalar(src->type)) {
2348 dest->def = src->def;
2349 } else {
2350 unsigned elems = glsl_get_length(src->type);
2351
2352 dest->elems = ralloc_array(mem_ctx, struct vtn_ssa_value *, elems);
2353 for (unsigned i = 0; i < elems; i++)
2354 dest->elems[i] = vtn_composite_copy(mem_ctx, src->elems[i]);
2355 }
2356
2357 return dest;
2358 }
2359
2360 static struct vtn_ssa_value *
vtn_composite_insert(struct vtn_builder * b,struct vtn_ssa_value * src,struct vtn_ssa_value * insert,const uint32_t * indices,unsigned num_indices)2361 vtn_composite_insert(struct vtn_builder *b, struct vtn_ssa_value *src,
2362 struct vtn_ssa_value *insert, const uint32_t *indices,
2363 unsigned num_indices)
2364 {
2365 struct vtn_ssa_value *dest = vtn_composite_copy(b, src);
2366
2367 struct vtn_ssa_value *cur = dest;
2368 unsigned i;
2369 for (i = 0; i < num_indices - 1; i++) {
2370 cur = cur->elems[indices[i]];
2371 }
2372
2373 if (glsl_type_is_vector_or_scalar(cur->type)) {
2374 /* According to the SPIR-V spec, OpCompositeInsert may work down to
2375 * the component granularity. In that case, the last index will be
2376 * the index to insert the scalar into the vector.
2377 */
2378
2379 cur->def = vtn_vector_insert(b, cur->def, insert->def, indices[i]);
2380 } else {
2381 cur->elems[indices[i]] = insert;
2382 }
2383
2384 return dest;
2385 }
2386
2387 static struct vtn_ssa_value *
vtn_composite_extract(struct vtn_builder * b,struct vtn_ssa_value * src,const uint32_t * indices,unsigned num_indices)2388 vtn_composite_extract(struct vtn_builder *b, struct vtn_ssa_value *src,
2389 const uint32_t *indices, unsigned num_indices)
2390 {
2391 struct vtn_ssa_value *cur = src;
2392 for (unsigned i = 0; i < num_indices; i++) {
2393 if (glsl_type_is_vector_or_scalar(cur->type)) {
2394 assert(i == num_indices - 1);
2395 /* According to the SPIR-V spec, OpCompositeExtract may work down to
2396 * the component granularity. The last index will be the index of the
2397 * vector to extract.
2398 */
2399
2400 struct vtn_ssa_value *ret = rzalloc(b, struct vtn_ssa_value);
2401 ret->type = glsl_scalar_type(glsl_get_base_type(cur->type));
2402 ret->def = vtn_vector_extract(b, cur->def, indices[i]);
2403 return ret;
2404 } else {
2405 cur = cur->elems[indices[i]];
2406 }
2407 }
2408
2409 return cur;
2410 }
2411
2412 static void
vtn_handle_composite(struct vtn_builder * b,SpvOp opcode,const uint32_t * w,unsigned count)2413 vtn_handle_composite(struct vtn_builder *b, SpvOp opcode,
2414 const uint32_t *w, unsigned count)
2415 {
2416 struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_ssa);
2417 const struct glsl_type *type =
2418 vtn_value(b, w[1], vtn_value_type_type)->type->type;
2419 val->ssa = vtn_create_ssa_value(b, type);
2420
2421 switch (opcode) {
2422 case SpvOpVectorExtractDynamic:
2423 val->ssa->def = vtn_vector_extract_dynamic(b, vtn_ssa_value(b, w[3])->def,
2424 vtn_ssa_value(b, w[4])->def);
2425 break;
2426
2427 case SpvOpVectorInsertDynamic:
2428 val->ssa->def = vtn_vector_insert_dynamic(b, vtn_ssa_value(b, w[3])->def,
2429 vtn_ssa_value(b, w[4])->def,
2430 vtn_ssa_value(b, w[5])->def);
2431 break;
2432
2433 case SpvOpVectorShuffle:
2434 val->ssa->def = vtn_vector_shuffle(b, glsl_get_vector_elements(type),
2435 vtn_ssa_value(b, w[3])->def,
2436 vtn_ssa_value(b, w[4])->def,
2437 w + 5);
2438 break;
2439
2440 case SpvOpCompositeConstruct: {
2441 unsigned elems = count - 3;
2442 if (glsl_type_is_vector_or_scalar(type)) {
2443 nir_ssa_def *srcs[4];
2444 for (unsigned i = 0; i < elems; i++)
2445 srcs[i] = vtn_ssa_value(b, w[3 + i])->def;
2446 val->ssa->def =
2447 vtn_vector_construct(b, glsl_get_vector_elements(type),
2448 elems, srcs);
2449 } else {
2450 val->ssa->elems = ralloc_array(b, struct vtn_ssa_value *, elems);
2451 for (unsigned i = 0; i < elems; i++)
2452 val->ssa->elems[i] = vtn_ssa_value(b, w[3 + i]);
2453 }
2454 break;
2455 }
2456 case SpvOpCompositeExtract:
2457 val->ssa = vtn_composite_extract(b, vtn_ssa_value(b, w[3]),
2458 w + 4, count - 4);
2459 break;
2460
2461 case SpvOpCompositeInsert:
2462 val->ssa = vtn_composite_insert(b, vtn_ssa_value(b, w[4]),
2463 vtn_ssa_value(b, w[3]),
2464 w + 5, count - 5);
2465 break;
2466
2467 case SpvOpCopyObject:
2468 val->ssa = vtn_composite_copy(b, vtn_ssa_value(b, w[3]));
2469 break;
2470
2471 default:
2472 unreachable("unknown composite operation");
2473 }
2474 }
2475
2476 static void
vtn_handle_barrier(struct vtn_builder * b,SpvOp opcode,const uint32_t * w,unsigned count)2477 vtn_handle_barrier(struct vtn_builder *b, SpvOp opcode,
2478 const uint32_t *w, unsigned count)
2479 {
2480 nir_intrinsic_op intrinsic_op;
2481 switch (opcode) {
2482 case SpvOpEmitVertex:
2483 case SpvOpEmitStreamVertex:
2484 intrinsic_op = nir_intrinsic_emit_vertex;
2485 break;
2486 case SpvOpEndPrimitive:
2487 case SpvOpEndStreamPrimitive:
2488 intrinsic_op = nir_intrinsic_end_primitive;
2489 break;
2490 case SpvOpMemoryBarrier:
2491 intrinsic_op = nir_intrinsic_memory_barrier;
2492 break;
2493 case SpvOpControlBarrier:
2494 intrinsic_op = nir_intrinsic_barrier;
2495 break;
2496 default:
2497 unreachable("unknown barrier instruction");
2498 }
2499
2500 nir_intrinsic_instr *intrin =
2501 nir_intrinsic_instr_create(b->shader, intrinsic_op);
2502
2503 if (opcode == SpvOpEmitStreamVertex || opcode == SpvOpEndStreamPrimitive)
2504 nir_intrinsic_set_stream_id(intrin, w[1]);
2505
2506 nir_builder_instr_insert(&b->nb, &intrin->instr);
2507 }
2508
2509 static unsigned
gl_primitive_from_spv_execution_mode(SpvExecutionMode mode)2510 gl_primitive_from_spv_execution_mode(SpvExecutionMode mode)
2511 {
2512 switch (mode) {
2513 case SpvExecutionModeInputPoints:
2514 case SpvExecutionModeOutputPoints:
2515 return 0; /* GL_POINTS */
2516 case SpvExecutionModeInputLines:
2517 return 1; /* GL_LINES */
2518 case SpvExecutionModeInputLinesAdjacency:
2519 return 0x000A; /* GL_LINE_STRIP_ADJACENCY_ARB */
2520 case SpvExecutionModeTriangles:
2521 return 4; /* GL_TRIANGLES */
2522 case SpvExecutionModeInputTrianglesAdjacency:
2523 return 0x000C; /* GL_TRIANGLES_ADJACENCY_ARB */
2524 case SpvExecutionModeQuads:
2525 return 7; /* GL_QUADS */
2526 case SpvExecutionModeIsolines:
2527 return 0x8E7A; /* GL_ISOLINES */
2528 case SpvExecutionModeOutputLineStrip:
2529 return 3; /* GL_LINE_STRIP */
2530 case SpvExecutionModeOutputTriangleStrip:
2531 return 5; /* GL_TRIANGLE_STRIP */
2532 default:
2533 assert(!"Invalid primitive type");
2534 return 4;
2535 }
2536 }
2537
2538 static unsigned
vertices_in_from_spv_execution_mode(SpvExecutionMode mode)2539 vertices_in_from_spv_execution_mode(SpvExecutionMode mode)
2540 {
2541 switch (mode) {
2542 case SpvExecutionModeInputPoints:
2543 return 1;
2544 case SpvExecutionModeInputLines:
2545 return 2;
2546 case SpvExecutionModeInputLinesAdjacency:
2547 return 4;
2548 case SpvExecutionModeTriangles:
2549 return 3;
2550 case SpvExecutionModeInputTrianglesAdjacency:
2551 return 6;
2552 default:
2553 assert(!"Invalid GS input mode");
2554 return 0;
2555 }
2556 }
2557
2558 static gl_shader_stage
stage_for_execution_model(SpvExecutionModel model)2559 stage_for_execution_model(SpvExecutionModel model)
2560 {
2561 switch (model) {
2562 case SpvExecutionModelVertex:
2563 return MESA_SHADER_VERTEX;
2564 case SpvExecutionModelTessellationControl:
2565 return MESA_SHADER_TESS_CTRL;
2566 case SpvExecutionModelTessellationEvaluation:
2567 return MESA_SHADER_TESS_EVAL;
2568 case SpvExecutionModelGeometry:
2569 return MESA_SHADER_GEOMETRY;
2570 case SpvExecutionModelFragment:
2571 return MESA_SHADER_FRAGMENT;
2572 case SpvExecutionModelGLCompute:
2573 return MESA_SHADER_COMPUTE;
2574 default:
2575 unreachable("Unsupported execution model");
2576 }
2577 }
2578
2579 #define spv_check_supported(name, cap) do { \
2580 if (!(b->ext && b->ext->name)) \
2581 vtn_warn("Unsupported SPIR-V capability: %s", \
2582 spirv_capability_to_string(cap)); \
2583 } while(0)
2584
2585 static bool
vtn_handle_preamble_instruction(struct vtn_builder * b,SpvOp opcode,const uint32_t * w,unsigned count)2586 vtn_handle_preamble_instruction(struct vtn_builder *b, SpvOp opcode,
2587 const uint32_t *w, unsigned count)
2588 {
2589 switch (opcode) {
2590 case SpvOpSource:
2591 case SpvOpSourceExtension:
2592 case SpvOpSourceContinued:
2593 case SpvOpExtension:
2594 /* Unhandled, but these are for debug so that's ok. */
2595 break;
2596
2597 case SpvOpCapability: {
2598 SpvCapability cap = w[1];
2599 switch (cap) {
2600 case SpvCapabilityMatrix:
2601 case SpvCapabilityShader:
2602 case SpvCapabilityGeometry:
2603 case SpvCapabilityGeometryPointSize:
2604 case SpvCapabilityUniformBufferArrayDynamicIndexing:
2605 case SpvCapabilitySampledImageArrayDynamicIndexing:
2606 case SpvCapabilityStorageBufferArrayDynamicIndexing:
2607 case SpvCapabilityStorageImageArrayDynamicIndexing:
2608 case SpvCapabilityImageRect:
2609 case SpvCapabilitySampledRect:
2610 case SpvCapabilitySampled1D:
2611 case SpvCapabilityImage1D:
2612 case SpvCapabilitySampledCubeArray:
2613 case SpvCapabilitySampledBuffer:
2614 case SpvCapabilityImageBuffer:
2615 case SpvCapabilityImageQuery:
2616 case SpvCapabilityDerivativeControl:
2617 case SpvCapabilityInterpolationFunction:
2618 case SpvCapabilityMultiViewport:
2619 case SpvCapabilitySampleRateShading:
2620 case SpvCapabilityClipDistance:
2621 case SpvCapabilityCullDistance:
2622 case SpvCapabilityInputAttachment:
2623 case SpvCapabilityImageGatherExtended:
2624 case SpvCapabilityStorageImageExtendedFormats:
2625 break;
2626
2627 case SpvCapabilityGeometryStreams:
2628 case SpvCapabilityLinkage:
2629 case SpvCapabilityVector16:
2630 case SpvCapabilityFloat16Buffer:
2631 case SpvCapabilityFloat16:
2632 case SpvCapabilityInt64:
2633 case SpvCapabilityInt64Atomics:
2634 case SpvCapabilityAtomicStorage:
2635 case SpvCapabilityInt16:
2636 case SpvCapabilityStorageImageMultisample:
2637 case SpvCapabilityImageCubeArray:
2638 case SpvCapabilityInt8:
2639 case SpvCapabilitySparseResidency:
2640 case SpvCapabilityMinLod:
2641 case SpvCapabilityTransformFeedback:
2642 case SpvCapabilityStorageImageReadWithoutFormat:
2643 case SpvCapabilityStorageImageWriteWithoutFormat:
2644 vtn_warn("Unsupported SPIR-V capability: %s",
2645 spirv_capability_to_string(cap));
2646 break;
2647
2648 case SpvCapabilityFloat64:
2649 spv_check_supported(float64, cap);
2650 break;
2651
2652 case SpvCapabilityAddresses:
2653 case SpvCapabilityKernel:
2654 case SpvCapabilityImageBasic:
2655 case SpvCapabilityImageReadWrite:
2656 case SpvCapabilityImageMipmap:
2657 case SpvCapabilityPipes:
2658 case SpvCapabilityGroups:
2659 case SpvCapabilityDeviceEnqueue:
2660 case SpvCapabilityLiteralSampler:
2661 case SpvCapabilityGenericPointer:
2662 vtn_warn("Unsupported OpenCL-style SPIR-V capability: %s",
2663 spirv_capability_to_string(cap));
2664 break;
2665
2666 case SpvCapabilityImageMSArray:
2667 spv_check_supported(image_ms_array, cap);
2668 break;
2669
2670 case SpvCapabilityTessellation:
2671 case SpvCapabilityTessellationPointSize:
2672 spv_check_supported(tessellation, cap);
2673 break;
2674 }
2675 break;
2676 }
2677
2678 case SpvOpExtInstImport:
2679 vtn_handle_extension(b, opcode, w, count);
2680 break;
2681
2682 case SpvOpMemoryModel:
2683 assert(w[1] == SpvAddressingModelLogical);
2684 assert(w[2] == SpvMemoryModelGLSL450);
2685 break;
2686
2687 case SpvOpEntryPoint: {
2688 struct vtn_value *entry_point = &b->values[w[2]];
2689 /* Let this be a name label regardless */
2690 unsigned name_words;
2691 entry_point->name = vtn_string_literal(b, &w[3], count - 3, &name_words);
2692
2693 if (strcmp(entry_point->name, b->entry_point_name) != 0 ||
2694 stage_for_execution_model(w[1]) != b->entry_point_stage)
2695 break;
2696
2697 assert(b->entry_point == NULL);
2698 b->entry_point = entry_point;
2699 break;
2700 }
2701
2702 case SpvOpString:
2703 vtn_push_value(b, w[1], vtn_value_type_string)->str =
2704 vtn_string_literal(b, &w[2], count - 2, NULL);
2705 break;
2706
2707 case SpvOpName:
2708 b->values[w[1]].name = vtn_string_literal(b, &w[2], count - 2, NULL);
2709 break;
2710
2711 case SpvOpMemberName:
2712 /* TODO */
2713 break;
2714
2715 case SpvOpExecutionMode:
2716 case SpvOpDecorationGroup:
2717 case SpvOpDecorate:
2718 case SpvOpMemberDecorate:
2719 case SpvOpGroupDecorate:
2720 case SpvOpGroupMemberDecorate:
2721 vtn_handle_decoration(b, opcode, w, count);
2722 break;
2723
2724 default:
2725 return false; /* End of preamble */
2726 }
2727
2728 return true;
2729 }
2730
2731 static void
vtn_handle_execution_mode(struct vtn_builder * b,struct vtn_value * entry_point,const struct vtn_decoration * mode,void * data)2732 vtn_handle_execution_mode(struct vtn_builder *b, struct vtn_value *entry_point,
2733 const struct vtn_decoration *mode, void *data)
2734 {
2735 assert(b->entry_point == entry_point);
2736
2737 switch(mode->exec_mode) {
2738 case SpvExecutionModeOriginUpperLeft:
2739 case SpvExecutionModeOriginLowerLeft:
2740 b->origin_upper_left =
2741 (mode->exec_mode == SpvExecutionModeOriginUpperLeft);
2742 break;
2743
2744 case SpvExecutionModeEarlyFragmentTests:
2745 assert(b->shader->stage == MESA_SHADER_FRAGMENT);
2746 b->shader->info->fs.early_fragment_tests = true;
2747 break;
2748
2749 case SpvExecutionModeInvocations:
2750 assert(b->shader->stage == MESA_SHADER_GEOMETRY);
2751 b->shader->info->gs.invocations = MAX2(1, mode->literals[0]);
2752 break;
2753
2754 case SpvExecutionModeDepthReplacing:
2755 assert(b->shader->stage == MESA_SHADER_FRAGMENT);
2756 b->shader->info->fs.depth_layout = FRAG_DEPTH_LAYOUT_ANY;
2757 break;
2758 case SpvExecutionModeDepthGreater:
2759 assert(b->shader->stage == MESA_SHADER_FRAGMENT);
2760 b->shader->info->fs.depth_layout = FRAG_DEPTH_LAYOUT_GREATER;
2761 break;
2762 case SpvExecutionModeDepthLess:
2763 assert(b->shader->stage == MESA_SHADER_FRAGMENT);
2764 b->shader->info->fs.depth_layout = FRAG_DEPTH_LAYOUT_LESS;
2765 break;
2766 case SpvExecutionModeDepthUnchanged:
2767 assert(b->shader->stage == MESA_SHADER_FRAGMENT);
2768 b->shader->info->fs.depth_layout = FRAG_DEPTH_LAYOUT_UNCHANGED;
2769 break;
2770
2771 case SpvExecutionModeLocalSize:
2772 assert(b->shader->stage == MESA_SHADER_COMPUTE);
2773 b->shader->info->cs.local_size[0] = mode->literals[0];
2774 b->shader->info->cs.local_size[1] = mode->literals[1];
2775 b->shader->info->cs.local_size[2] = mode->literals[2];
2776 break;
2777 case SpvExecutionModeLocalSizeHint:
2778 break; /* Nothing to do with this */
2779
2780 case SpvExecutionModeOutputVertices:
2781 if (b->shader->stage == MESA_SHADER_TESS_CTRL ||
2782 b->shader->stage == MESA_SHADER_TESS_EVAL) {
2783 b->shader->info->tess.tcs_vertices_out = mode->literals[0];
2784 } else {
2785 assert(b->shader->stage == MESA_SHADER_GEOMETRY);
2786 b->shader->info->gs.vertices_out = mode->literals[0];
2787 }
2788 break;
2789
2790 case SpvExecutionModeInputPoints:
2791 case SpvExecutionModeInputLines:
2792 case SpvExecutionModeInputLinesAdjacency:
2793 case SpvExecutionModeTriangles:
2794 case SpvExecutionModeInputTrianglesAdjacency:
2795 case SpvExecutionModeQuads:
2796 case SpvExecutionModeIsolines:
2797 if (b->shader->stage == MESA_SHADER_TESS_CTRL ||
2798 b->shader->stage == MESA_SHADER_TESS_EVAL) {
2799 b->shader->info->tess.primitive_mode =
2800 gl_primitive_from_spv_execution_mode(mode->exec_mode);
2801 } else {
2802 assert(b->shader->stage == MESA_SHADER_GEOMETRY);
2803 b->shader->info->gs.vertices_in =
2804 vertices_in_from_spv_execution_mode(mode->exec_mode);
2805 }
2806 break;
2807
2808 case SpvExecutionModeOutputPoints:
2809 case SpvExecutionModeOutputLineStrip:
2810 case SpvExecutionModeOutputTriangleStrip:
2811 assert(b->shader->stage == MESA_SHADER_GEOMETRY);
2812 b->shader->info->gs.output_primitive =
2813 gl_primitive_from_spv_execution_mode(mode->exec_mode);
2814 break;
2815
2816 case SpvExecutionModeSpacingEqual:
2817 assert(b->shader->stage == MESA_SHADER_TESS_CTRL ||
2818 b->shader->stage == MESA_SHADER_TESS_EVAL);
2819 b->shader->info->tess.spacing = TESS_SPACING_EQUAL;
2820 break;
2821 case SpvExecutionModeSpacingFractionalEven:
2822 assert(b->shader->stage == MESA_SHADER_TESS_CTRL ||
2823 b->shader->stage == MESA_SHADER_TESS_EVAL);
2824 b->shader->info->tess.spacing = TESS_SPACING_FRACTIONAL_EVEN;
2825 break;
2826 case SpvExecutionModeSpacingFractionalOdd:
2827 assert(b->shader->stage == MESA_SHADER_TESS_CTRL ||
2828 b->shader->stage == MESA_SHADER_TESS_EVAL);
2829 b->shader->info->tess.spacing = TESS_SPACING_FRACTIONAL_ODD;
2830 break;
2831 case SpvExecutionModeVertexOrderCw:
2832 assert(b->shader->stage == MESA_SHADER_TESS_CTRL ||
2833 b->shader->stage == MESA_SHADER_TESS_EVAL);
2834 /* Vulkan's notion of CCW seems to match the hardware backends,
2835 * but be the opposite of OpenGL. Currently NIR follows GL semantics,
2836 * so we set it backwards here.
2837 */
2838 b->shader->info->tess.ccw = true;
2839 break;
2840 case SpvExecutionModeVertexOrderCcw:
2841 assert(b->shader->stage == MESA_SHADER_TESS_CTRL ||
2842 b->shader->stage == MESA_SHADER_TESS_EVAL);
2843 /* Backwards; see above */
2844 b->shader->info->tess.ccw = false;
2845 break;
2846 case SpvExecutionModePointMode:
2847 assert(b->shader->stage == MESA_SHADER_TESS_CTRL ||
2848 b->shader->stage == MESA_SHADER_TESS_EVAL);
2849 b->shader->info->tess.point_mode = true;
2850 break;
2851
2852 case SpvExecutionModePixelCenterInteger:
2853 b->pixel_center_integer = true;
2854 break;
2855
2856 case SpvExecutionModeXfb:
2857 assert(!"Unhandled execution mode");
2858 break;
2859
2860 case SpvExecutionModeVecTypeHint:
2861 case SpvExecutionModeContractionOff:
2862 break; /* OpenCL */
2863 }
2864 }
2865
2866 static bool
vtn_handle_variable_or_type_instruction(struct vtn_builder * b,SpvOp opcode,const uint32_t * w,unsigned count)2867 vtn_handle_variable_or_type_instruction(struct vtn_builder *b, SpvOp opcode,
2868 const uint32_t *w, unsigned count)
2869 {
2870 switch (opcode) {
2871 case SpvOpSource:
2872 case SpvOpSourceContinued:
2873 case SpvOpSourceExtension:
2874 case SpvOpExtension:
2875 case SpvOpCapability:
2876 case SpvOpExtInstImport:
2877 case SpvOpMemoryModel:
2878 case SpvOpEntryPoint:
2879 case SpvOpExecutionMode:
2880 case SpvOpString:
2881 case SpvOpName:
2882 case SpvOpMemberName:
2883 case SpvOpDecorationGroup:
2884 case SpvOpDecorate:
2885 case SpvOpMemberDecorate:
2886 case SpvOpGroupDecorate:
2887 case SpvOpGroupMemberDecorate:
2888 assert(!"Invalid opcode types and variables section");
2889 break;
2890
2891 case SpvOpTypeVoid:
2892 case SpvOpTypeBool:
2893 case SpvOpTypeInt:
2894 case SpvOpTypeFloat:
2895 case SpvOpTypeVector:
2896 case SpvOpTypeMatrix:
2897 case SpvOpTypeImage:
2898 case SpvOpTypeSampler:
2899 case SpvOpTypeSampledImage:
2900 case SpvOpTypeArray:
2901 case SpvOpTypeRuntimeArray:
2902 case SpvOpTypeStruct:
2903 case SpvOpTypeOpaque:
2904 case SpvOpTypePointer:
2905 case SpvOpTypeFunction:
2906 case SpvOpTypeEvent:
2907 case SpvOpTypeDeviceEvent:
2908 case SpvOpTypeReserveId:
2909 case SpvOpTypeQueue:
2910 case SpvOpTypePipe:
2911 vtn_handle_type(b, opcode, w, count);
2912 break;
2913
2914 case SpvOpConstantTrue:
2915 case SpvOpConstantFalse:
2916 case SpvOpConstant:
2917 case SpvOpConstantComposite:
2918 case SpvOpConstantSampler:
2919 case SpvOpConstantNull:
2920 case SpvOpSpecConstantTrue:
2921 case SpvOpSpecConstantFalse:
2922 case SpvOpSpecConstant:
2923 case SpvOpSpecConstantComposite:
2924 case SpvOpSpecConstantOp:
2925 vtn_handle_constant(b, opcode, w, count);
2926 break;
2927
2928 case SpvOpUndef:
2929 case SpvOpVariable:
2930 vtn_handle_variables(b, opcode, w, count);
2931 break;
2932
2933 default:
2934 return false; /* End of preamble */
2935 }
2936
2937 return true;
2938 }
2939
2940 static bool
vtn_handle_body_instruction(struct vtn_builder * b,SpvOp opcode,const uint32_t * w,unsigned count)2941 vtn_handle_body_instruction(struct vtn_builder *b, SpvOp opcode,
2942 const uint32_t *w, unsigned count)
2943 {
2944 switch (opcode) {
2945 case SpvOpLabel:
2946 break;
2947
2948 case SpvOpLoopMerge:
2949 case SpvOpSelectionMerge:
2950 /* This is handled by cfg pre-pass and walk_blocks */
2951 break;
2952
2953 case SpvOpUndef: {
2954 struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_undef);
2955 val->type = vtn_value(b, w[1], vtn_value_type_type)->type;
2956 break;
2957 }
2958
2959 case SpvOpExtInst:
2960 vtn_handle_extension(b, opcode, w, count);
2961 break;
2962
2963 case SpvOpVariable:
2964 case SpvOpLoad:
2965 case SpvOpStore:
2966 case SpvOpCopyMemory:
2967 case SpvOpCopyMemorySized:
2968 case SpvOpAccessChain:
2969 case SpvOpInBoundsAccessChain:
2970 case SpvOpArrayLength:
2971 vtn_handle_variables(b, opcode, w, count);
2972 break;
2973
2974 case SpvOpFunctionCall:
2975 vtn_handle_function_call(b, opcode, w, count);
2976 break;
2977
2978 case SpvOpSampledImage:
2979 case SpvOpImage:
2980 case SpvOpImageSampleImplicitLod:
2981 case SpvOpImageSampleExplicitLod:
2982 case SpvOpImageSampleDrefImplicitLod:
2983 case SpvOpImageSampleDrefExplicitLod:
2984 case SpvOpImageSampleProjImplicitLod:
2985 case SpvOpImageSampleProjExplicitLod:
2986 case SpvOpImageSampleProjDrefImplicitLod:
2987 case SpvOpImageSampleProjDrefExplicitLod:
2988 case SpvOpImageFetch:
2989 case SpvOpImageGather:
2990 case SpvOpImageDrefGather:
2991 case SpvOpImageQuerySizeLod:
2992 case SpvOpImageQueryLod:
2993 case SpvOpImageQueryLevels:
2994 case SpvOpImageQuerySamples:
2995 vtn_handle_texture(b, opcode, w, count);
2996 break;
2997
2998 case SpvOpImageRead:
2999 case SpvOpImageWrite:
3000 case SpvOpImageTexelPointer:
3001 vtn_handle_image(b, opcode, w, count);
3002 break;
3003
3004 case SpvOpImageQuerySize: {
3005 struct vtn_access_chain *image =
3006 vtn_value(b, w[3], vtn_value_type_access_chain)->access_chain;
3007 if (glsl_type_is_image(image->var->var->interface_type)) {
3008 vtn_handle_image(b, opcode, w, count);
3009 } else {
3010 vtn_handle_texture(b, opcode, w, count);
3011 }
3012 break;
3013 }
3014
3015 case SpvOpAtomicLoad:
3016 case SpvOpAtomicExchange:
3017 case SpvOpAtomicCompareExchange:
3018 case SpvOpAtomicCompareExchangeWeak:
3019 case SpvOpAtomicIIncrement:
3020 case SpvOpAtomicIDecrement:
3021 case SpvOpAtomicIAdd:
3022 case SpvOpAtomicISub:
3023 case SpvOpAtomicSMin:
3024 case SpvOpAtomicUMin:
3025 case SpvOpAtomicSMax:
3026 case SpvOpAtomicUMax:
3027 case SpvOpAtomicAnd:
3028 case SpvOpAtomicOr:
3029 case SpvOpAtomicXor: {
3030 struct vtn_value *pointer = vtn_untyped_value(b, w[3]);
3031 if (pointer->value_type == vtn_value_type_image_pointer) {
3032 vtn_handle_image(b, opcode, w, count);
3033 } else {
3034 assert(pointer->value_type == vtn_value_type_access_chain);
3035 vtn_handle_ssbo_or_shared_atomic(b, opcode, w, count);
3036 }
3037 break;
3038 }
3039
3040 case SpvOpAtomicStore: {
3041 struct vtn_value *pointer = vtn_untyped_value(b, w[1]);
3042 if (pointer->value_type == vtn_value_type_image_pointer) {
3043 vtn_handle_image(b, opcode, w, count);
3044 } else {
3045 assert(pointer->value_type == vtn_value_type_access_chain);
3046 vtn_handle_ssbo_or_shared_atomic(b, opcode, w, count);
3047 }
3048 break;
3049 }
3050
3051 case SpvOpSNegate:
3052 case SpvOpFNegate:
3053 case SpvOpNot:
3054 case SpvOpAny:
3055 case SpvOpAll:
3056 case SpvOpConvertFToU:
3057 case SpvOpConvertFToS:
3058 case SpvOpConvertSToF:
3059 case SpvOpConvertUToF:
3060 case SpvOpUConvert:
3061 case SpvOpSConvert:
3062 case SpvOpFConvert:
3063 case SpvOpQuantizeToF16:
3064 case SpvOpConvertPtrToU:
3065 case SpvOpConvertUToPtr:
3066 case SpvOpPtrCastToGeneric:
3067 case SpvOpGenericCastToPtr:
3068 case SpvOpBitcast:
3069 case SpvOpIsNan:
3070 case SpvOpIsInf:
3071 case SpvOpIsFinite:
3072 case SpvOpIsNormal:
3073 case SpvOpSignBitSet:
3074 case SpvOpLessOrGreater:
3075 case SpvOpOrdered:
3076 case SpvOpUnordered:
3077 case SpvOpIAdd:
3078 case SpvOpFAdd:
3079 case SpvOpISub:
3080 case SpvOpFSub:
3081 case SpvOpIMul:
3082 case SpvOpFMul:
3083 case SpvOpUDiv:
3084 case SpvOpSDiv:
3085 case SpvOpFDiv:
3086 case SpvOpUMod:
3087 case SpvOpSRem:
3088 case SpvOpSMod:
3089 case SpvOpFRem:
3090 case SpvOpFMod:
3091 case SpvOpVectorTimesScalar:
3092 case SpvOpDot:
3093 case SpvOpIAddCarry:
3094 case SpvOpISubBorrow:
3095 case SpvOpUMulExtended:
3096 case SpvOpSMulExtended:
3097 case SpvOpShiftRightLogical:
3098 case SpvOpShiftRightArithmetic:
3099 case SpvOpShiftLeftLogical:
3100 case SpvOpLogicalEqual:
3101 case SpvOpLogicalNotEqual:
3102 case SpvOpLogicalOr:
3103 case SpvOpLogicalAnd:
3104 case SpvOpLogicalNot:
3105 case SpvOpBitwiseOr:
3106 case SpvOpBitwiseXor:
3107 case SpvOpBitwiseAnd:
3108 case SpvOpSelect:
3109 case SpvOpIEqual:
3110 case SpvOpFOrdEqual:
3111 case SpvOpFUnordEqual:
3112 case SpvOpINotEqual:
3113 case SpvOpFOrdNotEqual:
3114 case SpvOpFUnordNotEqual:
3115 case SpvOpULessThan:
3116 case SpvOpSLessThan:
3117 case SpvOpFOrdLessThan:
3118 case SpvOpFUnordLessThan:
3119 case SpvOpUGreaterThan:
3120 case SpvOpSGreaterThan:
3121 case SpvOpFOrdGreaterThan:
3122 case SpvOpFUnordGreaterThan:
3123 case SpvOpULessThanEqual:
3124 case SpvOpSLessThanEqual:
3125 case SpvOpFOrdLessThanEqual:
3126 case SpvOpFUnordLessThanEqual:
3127 case SpvOpUGreaterThanEqual:
3128 case SpvOpSGreaterThanEqual:
3129 case SpvOpFOrdGreaterThanEqual:
3130 case SpvOpFUnordGreaterThanEqual:
3131 case SpvOpDPdx:
3132 case SpvOpDPdy:
3133 case SpvOpFwidth:
3134 case SpvOpDPdxFine:
3135 case SpvOpDPdyFine:
3136 case SpvOpFwidthFine:
3137 case SpvOpDPdxCoarse:
3138 case SpvOpDPdyCoarse:
3139 case SpvOpFwidthCoarse:
3140 case SpvOpBitFieldInsert:
3141 case SpvOpBitFieldSExtract:
3142 case SpvOpBitFieldUExtract:
3143 case SpvOpBitReverse:
3144 case SpvOpBitCount:
3145 case SpvOpTranspose:
3146 case SpvOpOuterProduct:
3147 case SpvOpMatrixTimesScalar:
3148 case SpvOpVectorTimesMatrix:
3149 case SpvOpMatrixTimesVector:
3150 case SpvOpMatrixTimesMatrix:
3151 vtn_handle_alu(b, opcode, w, count);
3152 break;
3153
3154 case SpvOpVectorExtractDynamic:
3155 case SpvOpVectorInsertDynamic:
3156 case SpvOpVectorShuffle:
3157 case SpvOpCompositeConstruct:
3158 case SpvOpCompositeExtract:
3159 case SpvOpCompositeInsert:
3160 case SpvOpCopyObject:
3161 vtn_handle_composite(b, opcode, w, count);
3162 break;
3163
3164 case SpvOpEmitVertex:
3165 case SpvOpEndPrimitive:
3166 case SpvOpEmitStreamVertex:
3167 case SpvOpEndStreamPrimitive:
3168 case SpvOpControlBarrier:
3169 case SpvOpMemoryBarrier:
3170 vtn_handle_barrier(b, opcode, w, count);
3171 break;
3172
3173 default:
3174 unreachable("Unhandled opcode");
3175 }
3176
3177 return true;
3178 }
3179
3180 nir_function *
spirv_to_nir(const uint32_t * words,size_t word_count,struct nir_spirv_specialization * spec,unsigned num_spec,gl_shader_stage stage,const char * entry_point_name,const struct nir_spirv_supported_extensions * ext,const nir_shader_compiler_options * options)3181 spirv_to_nir(const uint32_t *words, size_t word_count,
3182 struct nir_spirv_specialization *spec, unsigned num_spec,
3183 gl_shader_stage stage, const char *entry_point_name,
3184 const struct nir_spirv_supported_extensions *ext,
3185 const nir_shader_compiler_options *options)
3186 {
3187 const uint32_t *word_end = words + word_count;
3188
3189 /* Handle the SPIR-V header (first 4 dwords) */
3190 assert(word_count > 5);
3191
3192 assert(words[0] == SpvMagicNumber);
3193 assert(words[1] >= 0x10000);
3194 /* words[2] == generator magic */
3195 unsigned value_id_bound = words[3];
3196 assert(words[4] == 0);
3197
3198 words+= 5;
3199
3200 /* Initialize the stn_builder object */
3201 struct vtn_builder *b = rzalloc(NULL, struct vtn_builder);
3202 b->value_id_bound = value_id_bound;
3203 b->values = rzalloc_array(b, struct vtn_value, value_id_bound);
3204 exec_list_make_empty(&b->functions);
3205 b->entry_point_stage = stage;
3206 b->entry_point_name = entry_point_name;
3207 b->ext = ext;
3208
3209 /* Handle all the preamble instructions */
3210 words = vtn_foreach_instruction(b, words, word_end,
3211 vtn_handle_preamble_instruction);
3212
3213 if (b->entry_point == NULL) {
3214 assert(!"Entry point not found");
3215 ralloc_free(b);
3216 return NULL;
3217 }
3218
3219 b->shader = nir_shader_create(NULL, stage, options, NULL);
3220
3221 /* Set shader info defaults */
3222 b->shader->info->gs.invocations = 1;
3223
3224 /* Parse execution modes */
3225 vtn_foreach_execution_mode(b, b->entry_point,
3226 vtn_handle_execution_mode, NULL);
3227
3228 b->specializations = spec;
3229 b->num_specializations = num_spec;
3230
3231 /* Handle all variable, type, and constant instructions */
3232 words = vtn_foreach_instruction(b, words, word_end,
3233 vtn_handle_variable_or_type_instruction);
3234
3235 vtn_build_cfg(b, words, word_end);
3236
3237 foreach_list_typed(struct vtn_function, func, node, &b->functions) {
3238 b->impl = func->impl;
3239 b->const_table = _mesa_hash_table_create(b, _mesa_hash_pointer,
3240 _mesa_key_pointer_equal);
3241
3242 vtn_function_emit(b, func, vtn_handle_body_instruction);
3243 }
3244
3245 assert(b->entry_point->value_type == vtn_value_type_function);
3246 nir_function *entry_point = b->entry_point->func->impl->function;
3247 assert(entry_point);
3248
3249 ralloc_free(b);
3250
3251 return entry_point;
3252 }
3253