src/compiler/spirv/spirv_to_nir.c

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