src/compiler/spirv/spirv_to_nir.c

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