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