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