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