src/compiler/spirv/spirv_to_nir.c

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