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