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