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