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