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    const enum glsl_sampler_dim sampler_dim = glsl_get_sampler_dim(image_type);
1339    const bool is_array = glsl_sampler_type_is_array(image_type);
1340    const bool is_shadow = glsl_sampler_type_is_shadow(image_type);
1341
1342    /* Figure out the base texture operation */
1343    nir_texop texop;
1344    switch (opcode) {
1345    case SpvOpImageSampleImplicitLod:
1346    case SpvOpImageSampleDrefImplicitLod:
1347    case SpvOpImageSampleProjImplicitLod:
1348    case SpvOpImageSampleProjDrefImplicitLod:
1349       texop = nir_texop_tex;
1350       break;
1351
1352    case SpvOpImageSampleExplicitLod:
1353    case SpvOpImageSampleDrefExplicitLod:
1354    case SpvOpImageSampleProjExplicitLod:
1355    case SpvOpImageSampleProjDrefExplicitLod:
1356       texop = nir_texop_txl;
1357       break;
1358
1359    case SpvOpImageFetch:
1360       if (glsl_get_sampler_dim(image_type) == GLSL_SAMPLER_DIM_MS) {
1361          texop = nir_texop_txf_ms;
1362       } else {
1363          texop = nir_texop_txf;
1364       }
1365       break;
1366
1367    case SpvOpImageGather:
1368    case SpvOpImageDrefGather:
1369       texop = nir_texop_tg4;
1370       break;
1371
1372    case SpvOpImageQuerySizeLod:
1373    case SpvOpImageQuerySize:
1374       texop = nir_texop_txs;
1375       break;
1376
1377    case SpvOpImageQueryLod:
1378       texop = nir_texop_lod;
1379       break;
1380
1381    case SpvOpImageQueryLevels:
1382       texop = nir_texop_query_levels;
1383       break;
1384
1385    case SpvOpImageQuerySamples:
1386       texop = nir_texop_texture_samples;
1387       break;
1388
1389    default:
1390       unreachable("Unhandled opcode");
1391    }
1392
1393    nir_tex_src srcs[8]; /* 8 should be enough */
1394    nir_tex_src *p = srcs;
1395
1396    unsigned idx = 4;
1397
1398    struct nir_ssa_def *coord;
1399    unsigned coord_components;
1400    switch (opcode) {
1401    case SpvOpImageSampleImplicitLod:
1402    case SpvOpImageSampleExplicitLod:
1403    case SpvOpImageSampleDrefImplicitLod:
1404    case SpvOpImageSampleDrefExplicitLod:
1405    case SpvOpImageSampleProjImplicitLod:
1406    case SpvOpImageSampleProjExplicitLod:
1407    case SpvOpImageSampleProjDrefImplicitLod:
1408    case SpvOpImageSampleProjDrefExplicitLod:
1409    case SpvOpImageFetch:
1410    case SpvOpImageGather:
1411    case SpvOpImageDrefGather:
1412    case SpvOpImageQueryLod: {
1413       /* All these types have the coordinate as their first real argument */
1414       switch (sampler_dim) {
1415       case GLSL_SAMPLER_DIM_1D:
1416       case GLSL_SAMPLER_DIM_BUF:
1417          coord_components = 1;
1418          break;
1419       case GLSL_SAMPLER_DIM_2D:
1420       case GLSL_SAMPLER_DIM_RECT:
1421       case GLSL_SAMPLER_DIM_MS:
1422          coord_components = 2;
1423          break;
1424       case GLSL_SAMPLER_DIM_3D:
1425       case GLSL_SAMPLER_DIM_CUBE:
1426          coord_components = 3;
1427          break;
1428       default:
1429          unreachable("Invalid sampler type");
1430       }
1431
1432       if (is_array && texop != nir_texop_lod)
1433          coord_components++;
1434
1435       coord = vtn_ssa_value(b, w[idx++])->def;
1436       p->src = nir_src_for_ssa(coord);
1437       p->src_type = nir_tex_src_coord;
1438       p++;
1439       break;
1440    }
1441
1442    default:
1443       coord = NULL;
1444       coord_components = 0;
1445       break;
1446    }
1447
1448    switch (opcode) {
1449    case SpvOpImageSampleProjImplicitLod:
1450    case SpvOpImageSampleProjExplicitLod:
1451    case SpvOpImageSampleProjDrefImplicitLod:
1452    case SpvOpImageSampleProjDrefExplicitLod:
1453       /* These have the projector as the last coordinate component */
1454       p->src = nir_src_for_ssa(nir_channel(&b->nb, coord, coord_components));
1455       p->src_type = nir_tex_src_projector;
1456       p++;
1457       break;
1458
1459    default:
1460       break;
1461    }
1462
1463    unsigned gather_component = 0;
1464    switch (opcode) {
1465    case SpvOpImageSampleDrefImplicitLod:
1466    case SpvOpImageSampleDrefExplicitLod:
1467    case SpvOpImageSampleProjDrefImplicitLod:
1468    case SpvOpImageSampleProjDrefExplicitLod:
1469    case SpvOpImageDrefGather:
1470       /* These all have an explicit depth value as their next source */
1471       (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_comparitor);
1472       break;
1473
1474    case SpvOpImageGather:
1475       /* This has a component as its next source */
1476       gather_component =
1477          vtn_value(b, w[idx++], vtn_value_type_constant)->constant->value.u[0];
1478       break;
1479
1480    default:
1481       break;
1482    }
1483
1484    /* For OpImageQuerySizeLod, we always have an LOD */
1485    if (opcode == SpvOpImageQuerySizeLod)
1486       (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_lod);
1487
1488    /* Now we need to handle some number of optional arguments */
1489    if (idx < count) {
1490       uint32_t operands = w[idx++];
1491
1492       if (operands & SpvImageOperandsBiasMask) {
1493          assert(texop == nir_texop_tex);
1494          texop = nir_texop_txb;
1495          (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_bias);
1496       }
1497
1498       if (operands & SpvImageOperandsLodMask) {
1499          assert(texop == nir_texop_txl || texop == nir_texop_txf ||
1500                 texop == nir_texop_txs);
1501          (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_lod);
1502       }
1503
1504       if (operands & SpvImageOperandsGradMask) {
1505          assert(texop == nir_texop_txl);
1506          texop = nir_texop_txd;
1507          (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_ddx);
1508          (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_ddy);
1509       }
1510
1511       if (operands & SpvImageOperandsOffsetMask ||
1512           operands & SpvImageOperandsConstOffsetMask)
1513          (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_offset);
1514
1515       if (operands & SpvImageOperandsConstOffsetsMask)
1516          assert(!"Constant offsets to texture gather not yet implemented");
1517
1518       if (operands & SpvImageOperandsSampleMask) {
1519          assert(texop == nir_texop_txf_ms);
1520          texop = nir_texop_txf_ms;
1521          (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_ms_index);
1522       }
1523    }
1524    /* We should have now consumed exactly all of the arguments */
1525    assert(idx == count);
1526
1527    nir_tex_instr *instr = nir_tex_instr_create(b->shader, p - srcs);
1528    instr->op = texop;
1529
1530    memcpy(instr->src, srcs, instr->num_srcs * sizeof(*instr->src));
1531
1532    instr->coord_components = coord_components;
1533    instr->sampler_dim = sampler_dim;
1534    instr->is_array = is_array;
1535    instr->is_shadow = is_shadow;
1536    instr->is_new_style_shadow =
1537       is_shadow && glsl_get_components(ret_type->type) == 1;
1538    instr->component = gather_component;
1539
1540    switch (glsl_get_sampler_result_type(image_type)) {
1541    case GLSL_TYPE_FLOAT:   instr->dest_type = nir_type_float;     break;
1542    case GLSL_TYPE_INT:     instr->dest_type = nir_type_int;       break;
1543    case GLSL_TYPE_UINT:    instr->dest_type = nir_type_uint;  break;
1544    case GLSL_TYPE_BOOL:    instr->dest_type = nir_type_bool;      break;
1545    default:
1546       unreachable("Invalid base type for sampler result");
1547    }
1548
1549    nir_deref_var *sampler = vtn_access_chain_to_deref(b, sampled.sampler);
1550    if (sampled.image) {
1551       nir_deref_var *image = vtn_access_chain_to_deref(b, sampled.image);
1552       instr->texture = nir_deref_as_var(nir_copy_deref(instr, &image->deref));
1553    } else {
1554       instr->texture = nir_deref_as_var(nir_copy_deref(instr, &sampler->deref));
1555    }
1556
1557    switch (instr->op) {
1558    case nir_texop_tex:
1559    case nir_texop_txb:
1560    case nir_texop_txl:
1561    case nir_texop_txd:
1562       /* These operations require a sampler */
1563       instr->sampler = nir_deref_as_var(nir_copy_deref(instr, &sampler->deref));
1564       break;
1565    case nir_texop_txf:
1566    case nir_texop_txf_ms:
1567    case nir_texop_txs:
1568    case nir_texop_lod:
1569    case nir_texop_tg4:
1570    case nir_texop_query_levels:
1571    case nir_texop_texture_samples:
1572    case nir_texop_samples_identical:
1573       /* These don't */
1574       instr->sampler = NULL;
1575       break;
1576    case nir_texop_txf_ms_mcs:
1577       unreachable("unexpected nir_texop_txf_ms_mcs");
1578    }
1579
1580    nir_ssa_dest_init(&instr->instr, &instr->dest,
1581                      nir_tex_instr_dest_size(instr), 32, NULL);
1582
1583    assert(glsl_get_vector_elements(ret_type->type) ==
1584           nir_tex_instr_dest_size(instr));
1585
1586    val->ssa = vtn_create_ssa_value(b, ret_type->type);
1587    val->ssa->def = &instr->dest.ssa;
1588
1589    nir_builder_instr_insert(&b->nb, &instr->instr);
1590 }
1591
1592 static void
1593 fill_common_atomic_sources(struct vtn_builder *b, SpvOp opcode,
1594                            const uint32_t *w, nir_src *src)
1595 {
1596    switch (opcode) {
1597    case SpvOpAtomicIIncrement:
1598       src[0] = nir_src_for_ssa(nir_imm_int(&b->nb, 1));
1599       break;
1600
1601    case SpvOpAtomicIDecrement:
1602       src[0] = nir_src_for_ssa(nir_imm_int(&b->nb, -1));
1603       break;
1604
1605    case SpvOpAtomicISub:
1606       src[0] =
1607          nir_src_for_ssa(nir_ineg(&b->nb, vtn_ssa_value(b, w[6])->def));
1608       break;
1609
1610    case SpvOpAtomicCompareExchange:
1611       src[0] = nir_src_for_ssa(vtn_ssa_value(b, w[8])->def);
1612       src[1] = nir_src_for_ssa(vtn_ssa_value(b, w[7])->def);
1613       break;
1614
1615    case SpvOpAtomicExchange:
1616    case SpvOpAtomicIAdd:
1617    case SpvOpAtomicSMin:
1618    case SpvOpAtomicUMin:
1619    case SpvOpAtomicSMax:
1620    case SpvOpAtomicUMax:
1621    case SpvOpAtomicAnd:
1622    case SpvOpAtomicOr:
1623    case SpvOpAtomicXor:
1624       src[0] = nir_src_for_ssa(vtn_ssa_value(b, w[6])->def);
1625       break;
1626
1627    default:
1628       unreachable("Invalid SPIR-V atomic");
1629    }
1630 }
1631
1632 static nir_ssa_def *
1633 get_image_coord(struct vtn_builder *b, uint32_t value)
1634 {
1635    struct vtn_ssa_value *coord = vtn_ssa_value(b, value);
1636
1637    /* The image_load_store intrinsics assume a 4-dim coordinate */
1638    unsigned dim = glsl_get_vector_elements(coord->type);
1639    unsigned swizzle[4];
1640    for (unsigned i = 0; i < 4; i++)
1641       swizzle[i] = MIN2(i, dim - 1);
1642
1643    return nir_swizzle(&b->nb, coord->def, swizzle, 4, false);
1644 }
1645
1646 static void
1647 vtn_handle_image(struct vtn_builder *b, SpvOp opcode,
1648                  const uint32_t *w, unsigned count)
1649 {
1650    /* Just get this one out of the way */
1651    if (opcode == SpvOpImageTexelPointer) {
1652       struct vtn_value *val =
1653          vtn_push_value(b, w[2], vtn_value_type_image_pointer);
1654       val->image = ralloc(b, struct vtn_image_pointer);
1655
1656       val->image->image =
1657          vtn_value(b, w[3], vtn_value_type_access_chain)->access_chain;
1658       val->image->coord = get_image_coord(b, w[4]);
1659       val->image->sample = vtn_ssa_value(b, w[5])->def;
1660       return;
1661    }
1662
1663    struct vtn_image_pointer image;
1664
1665    switch (opcode) {
1666    case SpvOpAtomicExchange:
1667    case SpvOpAtomicCompareExchange:
1668    case SpvOpAtomicCompareExchangeWeak:
1669    case SpvOpAtomicIIncrement:
1670    case SpvOpAtomicIDecrement:
1671    case SpvOpAtomicIAdd:
1672    case SpvOpAtomicISub:
1673    case SpvOpAtomicSMin:
1674    case SpvOpAtomicUMin:
1675    case SpvOpAtomicSMax:
1676    case SpvOpAtomicUMax:
1677    case SpvOpAtomicAnd:
1678    case SpvOpAtomicOr:
1679    case SpvOpAtomicXor:
1680       image = *vtn_value(b, w[3], vtn_value_type_image_pointer)->image;
1681       break;
1682
1683    case SpvOpAtomicLoad: {
1684       image.image =
1685          vtn_value(b, w[3], vtn_value_type_access_chain)->access_chain;
1686       break;
1687    }
1688
1689    case SpvOpAtomicStore: {
1690       image.image =
1691          vtn_value(b, w[1], vtn_value_type_access_chain)->access_chain;
1692       break;
1693    }
1694
1695    case SpvOpImageQuerySize:
1696       image.image =
1697          vtn_value(b, w[3], vtn_value_type_access_chain)->access_chain;
1698       image.coord = NULL;
1699       image.sample = NULL;
1700       break;
1701
1702    case SpvOpImageRead:
1703       image.image =
1704          vtn_value(b, w[3], vtn_value_type_access_chain)->access_chain;
1705       image.coord = get_image_coord(b, w[4]);
1706
1707       if (count > 5 && (w[5] & SpvImageOperandsSampleMask)) {
1708          assert(w[5] == SpvImageOperandsSampleMask);
1709          image.sample = vtn_ssa_value(b, w[6])->def;
1710       } else {
1711          image.sample = nir_ssa_undef(&b->nb, 1, 32);
1712       }
1713       break;
1714
1715    case SpvOpImageWrite:
1716       image.image =
1717          vtn_value(b, w[1], vtn_value_type_access_chain)->access_chain;
1718       image.coord = get_image_coord(b, w[2]);
1719
1720       /* texel = w[3] */
1721
1722       if (count > 4 && (w[4] & SpvImageOperandsSampleMask)) {
1723          assert(w[4] == SpvImageOperandsSampleMask);
1724          image.sample = vtn_ssa_value(b, w[5])->def;
1725       } else {
1726          image.sample = nir_ssa_undef(&b->nb, 1, 32);
1727       }
1728       break;
1729
1730    default:
1731       unreachable("Invalid image opcode");
1732    }
1733
1734    nir_intrinsic_op op;
1735    switch (opcode) {
1736 #define OP(S, N) case SpvOp##S: op = nir_intrinsic_image_##N; break;
1737    OP(ImageQuerySize,         size)
1738    OP(ImageRead,              load)
1739    OP(ImageWrite,             store)
1740    OP(AtomicLoad,             load)
1741    OP(AtomicStore,            store)
1742    OP(AtomicExchange,         atomic_exchange)
1743    OP(AtomicCompareExchange,  atomic_comp_swap)
1744    OP(AtomicIIncrement,       atomic_add)
1745    OP(AtomicIDecrement,       atomic_add)
1746    OP(AtomicIAdd,             atomic_add)
1747    OP(AtomicISub,             atomic_add)
1748    OP(AtomicSMin,             atomic_min)
1749    OP(AtomicUMin,             atomic_min)
1750    OP(AtomicSMax,             atomic_max)
1751    OP(AtomicUMax,             atomic_max)
1752    OP(AtomicAnd,              atomic_and)
1753    OP(AtomicOr,               atomic_or)
1754    OP(AtomicXor,              atomic_xor)
1755 #undef OP
1756    default:
1757       unreachable("Invalid image opcode");
1758    }
1759
1760    nir_intrinsic_instr *intrin = nir_intrinsic_instr_create(b->shader, op);
1761
1762    nir_deref_var *image_deref = vtn_access_chain_to_deref(b, image.image);
1763    intrin->variables[0] =
1764       nir_deref_as_var(nir_copy_deref(&intrin->instr, &image_deref->deref));
1765
1766    /* ImageQuerySize doesn't take any extra parameters */
1767    if (opcode != SpvOpImageQuerySize) {
1768       /* The image coordinate is always 4 components but we may not have that
1769        * many.  Swizzle to compensate.
1770        */
1771       unsigned swiz[4];
1772       for (unsigned i = 0; i < 4; i++)
1773          swiz[i] = i < image.coord->num_components ? i : 0;
1774       intrin->src[0] = nir_src_for_ssa(nir_swizzle(&b->nb, image.coord,
1775                                                    swiz, 4, false));
1776       intrin->src[1] = nir_src_for_ssa(image.sample);
1777    }
1778
1779    switch (opcode) {
1780    case SpvOpAtomicLoad:
1781    case SpvOpImageQuerySize:
1782    case SpvOpImageRead:
1783       break;
1784    case SpvOpAtomicStore:
1785       intrin->src[2] = nir_src_for_ssa(vtn_ssa_value(b, w[4])->def);
1786       break;
1787    case SpvOpImageWrite:
1788       intrin->src[2] = nir_src_for_ssa(vtn_ssa_value(b, w[3])->def);
1789       break;
1790
1791    case SpvOpAtomicIIncrement:
1792    case SpvOpAtomicIDecrement:
1793    case SpvOpAtomicExchange:
1794    case SpvOpAtomicIAdd:
1795    case SpvOpAtomicSMin:
1796    case SpvOpAtomicUMin:
1797    case SpvOpAtomicSMax:
1798    case SpvOpAtomicUMax:
1799    case SpvOpAtomicAnd:
1800    case SpvOpAtomicOr:
1801    case SpvOpAtomicXor:
1802       fill_common_atomic_sources(b, opcode, w, &intrin->src[2]);
1803       break;
1804
1805    default:
1806       unreachable("Invalid image opcode");
1807    }
1808
1809    if (opcode != SpvOpImageWrite) {
1810       struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_ssa);
1811       struct vtn_type *type = vtn_value(b, w[1], vtn_value_type_type)->type;
1812       nir_ssa_dest_init(&intrin->instr, &intrin->dest, 4, 32, NULL);
1813
1814       nir_builder_instr_insert(&b->nb, &intrin->instr);
1815
1816       /* The image intrinsics always return 4 channels but we may not want
1817        * that many.  Emit a mov to trim it down.
1818        */
1819       unsigned swiz[4] = {0, 1, 2, 3};
1820       val->ssa = vtn_create_ssa_value(b, type->type);
1821       val->ssa->def = nir_swizzle(&b->nb, &intrin->dest.ssa, swiz,
1822                                   glsl_get_vector_elements(type->type), false);
1823    } else {
1824       nir_builder_instr_insert(&b->nb, &intrin->instr);
1825    }
1826 }
1827
1828 static nir_intrinsic_op
1829 get_ssbo_nir_atomic_op(SpvOp opcode)
1830 {
1831    switch (opcode) {
1832    case SpvOpAtomicLoad:      return nir_intrinsic_load_ssbo;
1833    case SpvOpAtomicStore:     return nir_intrinsic_store_ssbo;
1834 #define OP(S, N) case SpvOp##S: return nir_intrinsic_ssbo_##N;
1835    OP(AtomicExchange,         atomic_exchange)
1836    OP(AtomicCompareExchange,  atomic_comp_swap)
1837    OP(AtomicIIncrement,       atomic_add)
1838    OP(AtomicIDecrement,       atomic_add)
1839    OP(AtomicIAdd,             atomic_add)
1840    OP(AtomicISub,             atomic_add)
1841    OP(AtomicSMin,             atomic_imin)
1842    OP(AtomicUMin,             atomic_umin)
1843    OP(AtomicSMax,             atomic_imax)
1844    OP(AtomicUMax,             atomic_umax)
1845    OP(AtomicAnd,              atomic_and)
1846    OP(AtomicOr,               atomic_or)
1847    OP(AtomicXor,              atomic_xor)
1848 #undef OP
1849    default:
1850       unreachable("Invalid SSBO atomic");
1851    }
1852 }
1853
1854 static nir_intrinsic_op
1855 get_shared_nir_atomic_op(SpvOp opcode)
1856 {
1857    switch (opcode) {
1858    case SpvOpAtomicLoad:      return nir_intrinsic_load_var;
1859    case SpvOpAtomicStore:     return nir_intrinsic_store_var;
1860 #define OP(S, N) case SpvOp##S: return nir_intrinsic_var_##N;
1861    OP(AtomicExchange,         atomic_exchange)
1862    OP(AtomicCompareExchange,  atomic_comp_swap)
1863    OP(AtomicIIncrement,       atomic_add)
1864    OP(AtomicIDecrement,       atomic_add)
1865    OP(AtomicIAdd,             atomic_add)
1866    OP(AtomicISub,             atomic_add)
1867    OP(AtomicSMin,             atomic_imin)
1868    OP(AtomicUMin,             atomic_umin)
1869    OP(AtomicSMax,             atomic_imax)
1870    OP(AtomicUMax,             atomic_umax)
1871    OP(AtomicAnd,              atomic_and)
1872    OP(AtomicOr,               atomic_or)
1873    OP(AtomicXor,              atomic_xor)
1874 #undef OP
1875    default:
1876       unreachable("Invalid shared atomic");
1877    }
1878 }
1879
1880 static void
1881 vtn_handle_ssbo_or_shared_atomic(struct vtn_builder *b, SpvOp opcode,
1882                                  const uint32_t *w, unsigned count)
1883 {
1884    struct vtn_access_chain *chain;
1885    nir_intrinsic_instr *atomic;
1886
1887    switch (opcode) {
1888    case SpvOpAtomicLoad:
1889    case SpvOpAtomicExchange:
1890    case SpvOpAtomicCompareExchange:
1891    case SpvOpAtomicCompareExchangeWeak:
1892    case SpvOpAtomicIIncrement:
1893    case SpvOpAtomicIDecrement:
1894    case SpvOpAtomicIAdd:
1895    case SpvOpAtomicISub:
1896    case SpvOpAtomicSMin:
1897    case SpvOpAtomicUMin:
1898    case SpvOpAtomicSMax:
1899    case SpvOpAtomicUMax:
1900    case SpvOpAtomicAnd:
1901    case SpvOpAtomicOr:
1902    case SpvOpAtomicXor:
1903       chain =
1904          vtn_value(b, w[3], vtn_value_type_access_chain)->access_chain;
1905       break;
1906
1907    case SpvOpAtomicStore:
1908       chain =
1909          vtn_value(b, w[1], vtn_value_type_access_chain)->access_chain;
1910       break;
1911
1912    default:
1913       unreachable("Invalid SPIR-V atomic");
1914    }
1915
1916    /*
1917    SpvScope scope = w[4];
1918    SpvMemorySemanticsMask semantics = w[5];
1919    */
1920
1921    if (chain->var->mode == vtn_variable_mode_workgroup) {
1922       nir_deref *deref = &vtn_access_chain_to_deref(b, chain)->deref;
1923       nir_intrinsic_op op = get_shared_nir_atomic_op(opcode);
1924       atomic = nir_intrinsic_instr_create(b->nb.shader, op);
1925       atomic->variables[0] = nir_deref_as_var(nir_copy_deref(atomic, deref));
1926       fill_common_atomic_sources(b, opcode, w, &atomic->src[0]);
1927    } else {
1928       assert(chain->var->mode == vtn_variable_mode_ssbo);
1929       struct vtn_type *type;
1930       nir_ssa_def *offset, *index;
1931       offset = vtn_access_chain_to_offset(b, chain, &index, &type, NULL, false);
1932
1933       nir_intrinsic_op op = get_ssbo_nir_atomic_op(opcode);
1934
1935       atomic = nir_intrinsic_instr_create(b->nb.shader, op);
1936
1937       switch (opcode) {
1938       case SpvOpAtomicLoad:
1939          atomic->num_components = glsl_get_vector_elements(type->type);
1940          atomic->src[0] = nir_src_for_ssa(index);
1941          atomic->src[1] = nir_src_for_ssa(offset);
1942          break;
1943
1944       case SpvOpAtomicStore:
1945          atomic->num_components = glsl_get_vector_elements(type->type);
1946          nir_intrinsic_set_write_mask(atomic, (1 << atomic->num_components) - 1);
1947          atomic->src[0] = nir_src_for_ssa(vtn_ssa_value(b, w[4])->def);
1948          atomic->src[1] = nir_src_for_ssa(index);
1949          atomic->src[2] = nir_src_for_ssa(offset);
1950          break;
1951
1952       case SpvOpAtomicExchange:
1953       case SpvOpAtomicCompareExchange:
1954       case SpvOpAtomicCompareExchangeWeak:
1955       case SpvOpAtomicIIncrement:
1956       case SpvOpAtomicIDecrement:
1957       case SpvOpAtomicIAdd:
1958       case SpvOpAtomicISub:
1959       case SpvOpAtomicSMin:
1960       case SpvOpAtomicUMin:
1961       case SpvOpAtomicSMax:
1962       case SpvOpAtomicUMax:
1963       case SpvOpAtomicAnd:
1964       case SpvOpAtomicOr:
1965       case SpvOpAtomicXor:
1966          atomic->src[0] = nir_src_for_ssa(index);
1967          atomic->src[1] = nir_src_for_ssa(offset);
1968          fill_common_atomic_sources(b, opcode, w, &atomic->src[2]);
1969          break;
1970
1971       default:
1972          unreachable("Invalid SPIR-V atomic");
1973       }
1974    }
1975
1976    if (opcode != SpvOpAtomicStore) {
1977       struct vtn_type *type = vtn_value(b, w[1], vtn_value_type_type)->type;
1978
1979       nir_ssa_dest_init(&atomic->instr, &atomic->dest,
1980                         glsl_get_vector_elements(type->type),
1981                         glsl_get_bit_size(type->type), NULL);
1982
1983       struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_ssa);
1984       val->ssa = rzalloc(b, struct vtn_ssa_value);
1985       val->ssa->def = &atomic->dest.ssa;
1986       val->ssa->type = type->type;
1987    }
1988
1989    nir_builder_instr_insert(&b->nb, &atomic->instr);
1990 }
1991
1992 static nir_alu_instr *
1993 create_vec(nir_shader *shader, unsigned num_components, unsigned bit_size)
1994 {
1995    nir_op op;
1996    switch (num_components) {
1997    case 1: op = nir_op_fmov; break;
1998    case 2: op = nir_op_vec2; break;
1999    case 3: op = nir_op_vec3; break;
2000    case 4: op = nir_op_vec4; break;
2001    default: unreachable("bad vector size");
2002    }
2003
2004    nir_alu_instr *vec = nir_alu_instr_create(shader, op);
2005    nir_ssa_dest_init(&vec->instr, &vec->dest.dest, num_components,
2006                      bit_size, NULL);
2007    vec->dest.write_mask = (1 << num_components) - 1;
2008
2009    return vec;
2010 }
2011
2012 struct vtn_ssa_value *
2013 vtn_ssa_transpose(struct vtn_builder *b, struct vtn_ssa_value *src)
2014 {
2015    if (src->transposed)
2016       return src->transposed;
2017
2018    struct vtn_ssa_value *dest =
2019       vtn_create_ssa_value(b, glsl_transposed_type(src->type));
2020
2021    for (unsigned i = 0; i < glsl_get_matrix_columns(dest->type); i++) {
2022       nir_alu_instr *vec = create_vec(b->shader,
2023                                       glsl_get_matrix_columns(src->type),
2024                                       glsl_get_bit_size(src->type));
2025       if (glsl_type_is_vector_or_scalar(src->type)) {
2026           vec->src[0].src = nir_src_for_ssa(src->def);
2027           vec->src[0].swizzle[0] = i;
2028       } else {
2029          for (unsigned j = 0; j < glsl_get_matrix_columns(src->type); j++) {
2030             vec->src[j].src = nir_src_for_ssa(src->elems[j]->def);
2031             vec->src[j].swizzle[0] = i;
2032          }
2033       }
2034       nir_builder_instr_insert(&b->nb, &vec->instr);
2035       dest->elems[i]->def = &vec->dest.dest.ssa;
2036    }
2037
2038    dest->transposed = src;
2039
2040    return dest;
2041 }
2042
2043 nir_ssa_def *
2044 vtn_vector_extract(struct vtn_builder *b, nir_ssa_def *src, unsigned index)
2045 {
2046    unsigned swiz[4] = { index };
2047    return nir_swizzle(&b->nb, src, swiz, 1, true);
2048 }
2049
2050 nir_ssa_def *
2051 vtn_vector_insert(struct vtn_builder *b, nir_ssa_def *src, nir_ssa_def *insert,
2052                   unsigned index)
2053 {
2054    nir_alu_instr *vec = create_vec(b->shader, src->num_components,
2055                                    src->bit_size);
2056
2057    for (unsigned i = 0; i < src->num_components; i++) {
2058       if (i == index) {
2059          vec->src[i].src = nir_src_for_ssa(insert);
2060       } else {
2061          vec->src[i].src = nir_src_for_ssa(src);
2062          vec->src[i].swizzle[0] = i;
2063       }
2064    }
2065
2066    nir_builder_instr_insert(&b->nb, &vec->instr);
2067
2068    return &vec->dest.dest.ssa;
2069 }
2070
2071 nir_ssa_def *
2072 vtn_vector_extract_dynamic(struct vtn_builder *b, nir_ssa_def *src,
2073                            nir_ssa_def *index)
2074 {
2075    nir_ssa_def *dest = vtn_vector_extract(b, src, 0);
2076    for (unsigned i = 1; i < src->num_components; i++)
2077       dest = nir_bcsel(&b->nb, nir_ieq(&b->nb, index, nir_imm_int(&b->nb, i)),
2078                        vtn_vector_extract(b, src, i), dest);
2079
2080    return dest;
2081 }
2082
2083 nir_ssa_def *
2084 vtn_vector_insert_dynamic(struct vtn_builder *b, nir_ssa_def *src,
2085                           nir_ssa_def *insert, nir_ssa_def *index)
2086 {
2087    nir_ssa_def *dest = vtn_vector_insert(b, src, insert, 0);
2088    for (unsigned i = 1; i < src->num_components; i++)
2089       dest = nir_bcsel(&b->nb, nir_ieq(&b->nb, index, nir_imm_int(&b->nb, i)),
2090                        vtn_vector_insert(b, src, insert, i), dest);
2091
2092    return dest;
2093 }
2094
2095 static nir_ssa_def *
2096 vtn_vector_shuffle(struct vtn_builder *b, unsigned num_components,
2097                    nir_ssa_def *src0, nir_ssa_def *src1,
2098                    const uint32_t *indices)
2099 {
2100    nir_alu_instr *vec = create_vec(b->shader, num_components, src0->bit_size);
2101
2102    for (unsigned i = 0; i < num_components; i++) {
2103       uint32_t index = indices[i];
2104       if (index == 0xffffffff) {
2105          vec->src[i].src =
2106             nir_src_for_ssa(nir_ssa_undef(&b->nb, 1, src0->bit_size));
2107       } else if (index < src0->num_components) {
2108          vec->src[i].src = nir_src_for_ssa(src0);
2109          vec->src[i].swizzle[0] = index;
2110       } else {
2111          vec->src[i].src = nir_src_for_ssa(src1);
2112          vec->src[i].swizzle[0] = index - src0->num_components;
2113       }
2114    }
2115
2116    nir_builder_instr_insert(&b->nb, &vec->instr);
2117
2118    return &vec->dest.dest.ssa;
2119 }
2120
2121 /*
2122  * Concatentates a number of vectors/scalars together to produce a vector
2123  */
2124 static nir_ssa_def *
2125 vtn_vector_construct(struct vtn_builder *b, unsigned num_components,
2126                      unsigned num_srcs, nir_ssa_def **srcs)
2127 {
2128    nir_alu_instr *vec = create_vec(b->shader, num_components,
2129                                    srcs[0]->bit_size);
2130
2131    unsigned dest_idx = 0;
2132    for (unsigned i = 0; i < num_srcs; i++) {
2133       nir_ssa_def *src = srcs[i];
2134       for (unsigned j = 0; j < src->num_components; j++) {
2135          vec->src[dest_idx].src = nir_src_for_ssa(src);
2136          vec->src[dest_idx].swizzle[0] = j;
2137          dest_idx++;
2138       }
2139    }
2140
2141    nir_builder_instr_insert(&b->nb, &vec->instr);
2142
2143    return &vec->dest.dest.ssa;
2144 }
2145
2146 static struct vtn_ssa_value *
2147 vtn_composite_copy(void *mem_ctx, struct vtn_ssa_value *src)
2148 {
2149    struct vtn_ssa_value *dest = rzalloc(mem_ctx, struct vtn_ssa_value);
2150    dest->type = src->type;
2151
2152    if (glsl_type_is_vector_or_scalar(src->type)) {
2153       dest->def = src->def;
2154    } else {
2155       unsigned elems = glsl_get_length(src->type);
2156
2157       dest->elems = ralloc_array(mem_ctx, struct vtn_ssa_value *, elems);
2158       for (unsigned i = 0; i < elems; i++)
2159          dest->elems[i] = vtn_composite_copy(mem_ctx, src->elems[i]);
2160    }
2161
2162    return dest;
2163 }
2164
2165 static struct vtn_ssa_value *
2166 vtn_composite_insert(struct vtn_builder *b, struct vtn_ssa_value *src,
2167                      struct vtn_ssa_value *insert, const uint32_t *indices,
2168                      unsigned num_indices)
2169 {
2170    struct vtn_ssa_value *dest = vtn_composite_copy(b, src);
2171
2172    struct vtn_ssa_value *cur = dest;
2173    unsigned i;
2174    for (i = 0; i < num_indices - 1; i++) {
2175       cur = cur->elems[indices[i]];
2176    }
2177
2178    if (glsl_type_is_vector_or_scalar(cur->type)) {
2179       /* According to the SPIR-V spec, OpCompositeInsert may work down to
2180        * the component granularity. In that case, the last index will be
2181        * the index to insert the scalar into the vector.
2182        */
2183
2184       cur->def = vtn_vector_insert(b, cur->def, insert->def, indices[i]);
2185    } else {
2186       cur->elems[indices[i]] = insert;
2187    }
2188
2189    return dest;
2190 }
2191
2192 static struct vtn_ssa_value *
2193 vtn_composite_extract(struct vtn_builder *b, struct vtn_ssa_value *src,
2194                       const uint32_t *indices, unsigned num_indices)
2195 {
2196    struct vtn_ssa_value *cur = src;
2197    for (unsigned i = 0; i < num_indices; i++) {
2198       if (glsl_type_is_vector_or_scalar(cur->type)) {
2199          assert(i == num_indices - 1);
2200          /* According to the SPIR-V spec, OpCompositeExtract may work down to
2201           * the component granularity. The last index will be the index of the
2202           * vector to extract.
2203           */
2204
2205          struct vtn_ssa_value *ret = rzalloc(b, struct vtn_ssa_value);
2206          ret->type = glsl_scalar_type(glsl_get_base_type(cur->type));
2207          ret->def = vtn_vector_extract(b, cur->def, indices[i]);
2208          return ret;
2209       } else {
2210          cur = cur->elems[indices[i]];
2211       }
2212    }
2213
2214    return cur;
2215 }
2216
2217 static void
2218 vtn_handle_composite(struct vtn_builder *b, SpvOp opcode,
2219                      const uint32_t *w, unsigned count)
2220 {
2221    struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_ssa);
2222    const struct glsl_type *type =
2223       vtn_value(b, w[1], vtn_value_type_type)->type->type;
2224    val->ssa = vtn_create_ssa_value(b, type);
2225
2226    switch (opcode) {
2227    case SpvOpVectorExtractDynamic:
2228       val->ssa->def = vtn_vector_extract_dynamic(b, vtn_ssa_value(b, w[3])->def,
2229                                                  vtn_ssa_value(b, w[4])->def);
2230       break;
2231
2232    case SpvOpVectorInsertDynamic:
2233       val->ssa->def = vtn_vector_insert_dynamic(b, vtn_ssa_value(b, w[3])->def,
2234                                                 vtn_ssa_value(b, w[4])->def,
2235                                                 vtn_ssa_value(b, w[5])->def);
2236       break;
2237
2238    case SpvOpVectorShuffle:
2239       val->ssa->def = vtn_vector_shuffle(b, glsl_get_vector_elements(type),
2240                                          vtn_ssa_value(b, w[3])->def,
2241                                          vtn_ssa_value(b, w[4])->def,
2242                                          w + 5);
2243       break;
2244
2245    case SpvOpCompositeConstruct: {
2246       unsigned elems = count - 3;
2247       if (glsl_type_is_vector_or_scalar(type)) {
2248          nir_ssa_def *srcs[4];
2249          for (unsigned i = 0; i < elems; i++)
2250             srcs[i] = vtn_ssa_value(b, w[3 + i])->def;
2251          val->ssa->def =
2252             vtn_vector_construct(b, glsl_get_vector_elements(type),
2253                                  elems, srcs);
2254       } else {
2255          val->ssa->elems = ralloc_array(b, struct vtn_ssa_value *, elems);
2256          for (unsigned i = 0; i < elems; i++)
2257             val->ssa->elems[i] = vtn_ssa_value(b, w[3 + i]);
2258       }
2259       break;
2260    }
2261    case SpvOpCompositeExtract:
2262       val->ssa = vtn_composite_extract(b, vtn_ssa_value(b, w[3]),
2263                                        w + 4, count - 4);
2264       break;
2265
2266    case SpvOpCompositeInsert:
2267       val->ssa = vtn_composite_insert(b, vtn_ssa_value(b, w[4]),
2268                                       vtn_ssa_value(b, w[3]),
2269                                       w + 5, count - 5);
2270       break;
2271
2272    case SpvOpCopyObject:
2273       val->ssa = vtn_composite_copy(b, vtn_ssa_value(b, w[3]));
2274       break;
2275
2276    default:
2277       unreachable("unknown composite operation");
2278    }
2279 }
2280
2281 static void
2282 vtn_handle_barrier(struct vtn_builder *b, SpvOp opcode,
2283                    const uint32_t *w, unsigned count)
2284 {
2285    nir_intrinsic_op intrinsic_op;
2286    switch (opcode) {
2287    case SpvOpEmitVertex:
2288    case SpvOpEmitStreamVertex:
2289       intrinsic_op = nir_intrinsic_emit_vertex;
2290       break;
2291    case SpvOpEndPrimitive:
2292    case SpvOpEndStreamPrimitive:
2293       intrinsic_op = nir_intrinsic_end_primitive;
2294       break;
2295    case SpvOpMemoryBarrier:
2296       intrinsic_op = nir_intrinsic_memory_barrier;
2297       break;
2298    case SpvOpControlBarrier:
2299       intrinsic_op = nir_intrinsic_barrier;
2300       break;
2301    default:
2302       unreachable("unknown barrier instruction");
2303    }
2304
2305    nir_intrinsic_instr *intrin =
2306       nir_intrinsic_instr_create(b->shader, intrinsic_op);
2307
2308    if (opcode == SpvOpEmitStreamVertex || opcode == SpvOpEndStreamPrimitive)
2309       nir_intrinsic_set_stream_id(intrin, w[1]);
2310
2311    nir_builder_instr_insert(&b->nb, &intrin->instr);
2312 }
2313
2314 static unsigned
2315 gl_primitive_from_spv_execution_mode(SpvExecutionMode mode)
2316 {
2317    switch (mode) {
2318    case SpvExecutionModeInputPoints:
2319    case SpvExecutionModeOutputPoints:
2320       return 0; /* GL_POINTS */
2321    case SpvExecutionModeInputLines:
2322       return 1; /* GL_LINES */
2323    case SpvExecutionModeInputLinesAdjacency:
2324       return 0x000A; /* GL_LINE_STRIP_ADJACENCY_ARB */
2325    case SpvExecutionModeTriangles:
2326       return 4; /* GL_TRIANGLES */
2327    case SpvExecutionModeInputTrianglesAdjacency:
2328       return 0x000C; /* GL_TRIANGLES_ADJACENCY_ARB */
2329    case SpvExecutionModeQuads:
2330       return 7; /* GL_QUADS */
2331    case SpvExecutionModeIsolines:
2332       return 0x8E7A; /* GL_ISOLINES */
2333    case SpvExecutionModeOutputLineStrip:
2334       return 3; /* GL_LINE_STRIP */
2335    case SpvExecutionModeOutputTriangleStrip:
2336       return 5; /* GL_TRIANGLE_STRIP */
2337    default:
2338       assert(!"Invalid primitive type");
2339       return 4;
2340    }
2341 }
2342
2343 static unsigned
2344 vertices_in_from_spv_execution_mode(SpvExecutionMode mode)
2345 {
2346    switch (mode) {
2347    case SpvExecutionModeInputPoints:
2348       return 1;
2349    case SpvExecutionModeInputLines:
2350       return 2;
2351    case SpvExecutionModeInputLinesAdjacency:
2352       return 4;
2353    case SpvExecutionModeTriangles:
2354       return 3;
2355    case SpvExecutionModeInputTrianglesAdjacency:
2356       return 6;
2357    default:
2358       assert(!"Invalid GS input mode");
2359       return 0;
2360    }
2361 }
2362
2363 static gl_shader_stage
2364 stage_for_execution_model(SpvExecutionModel model)
2365 {
2366    switch (model) {
2367    case SpvExecutionModelVertex:
2368       return MESA_SHADER_VERTEX;
2369    case SpvExecutionModelTessellationControl:
2370       return MESA_SHADER_TESS_CTRL;
2371    case SpvExecutionModelTessellationEvaluation:
2372       return MESA_SHADER_TESS_EVAL;
2373    case SpvExecutionModelGeometry:
2374       return MESA_SHADER_GEOMETRY;
2375    case SpvExecutionModelFragment:
2376       return MESA_SHADER_FRAGMENT;
2377    case SpvExecutionModelGLCompute:
2378       return MESA_SHADER_COMPUTE;
2379    default:
2380       unreachable("Unsupported execution model");
2381    }
2382 }
2383
2384 static bool
2385 vtn_handle_preamble_instruction(struct vtn_builder *b, SpvOp opcode,
2386                                 const uint32_t *w, unsigned count)
2387 {
2388    switch (opcode) {
2389    case SpvOpSource:
2390    case SpvOpSourceExtension:
2391    case SpvOpSourceContinued:
2392    case SpvOpExtension:
2393       /* Unhandled, but these are for debug so that's ok. */
2394       break;
2395
2396    case SpvOpCapability: {
2397       SpvCapability cap = w[1];
2398       switch (cap) {
2399       case SpvCapabilityMatrix:
2400       case SpvCapabilityShader:
2401       case SpvCapabilityGeometry:
2402       case SpvCapabilityGeometryPointSize:
2403       case SpvCapabilityUniformBufferArrayDynamicIndexing:
2404       case SpvCapabilitySampledImageArrayDynamicIndexing:
2405       case SpvCapabilityStorageBufferArrayDynamicIndexing:
2406       case SpvCapabilityStorageImageArrayDynamicIndexing:
2407       case SpvCapabilityImageRect:
2408       case SpvCapabilitySampledRect:
2409       case SpvCapabilitySampled1D:
2410       case SpvCapabilityImage1D:
2411       case SpvCapabilitySampledCubeArray:
2412       case SpvCapabilitySampledBuffer:
2413       case SpvCapabilityImageBuffer:
2414       case SpvCapabilityImageQuery:
2415       case SpvCapabilityDerivativeControl:
2416       case SpvCapabilityInterpolationFunction:
2417       case SpvCapabilityMultiViewport:
2418          break;
2419
2420       case SpvCapabilityClipDistance:
2421       case SpvCapabilityCullDistance:
2422       case SpvCapabilityGeometryStreams:
2423       case SpvCapabilityTessellation:
2424       case SpvCapabilityTessellationPointSize:
2425       case SpvCapabilityLinkage:
2426       case SpvCapabilityVector16:
2427       case SpvCapabilityFloat16Buffer:
2428       case SpvCapabilityFloat16:
2429       case SpvCapabilityFloat64:
2430       case SpvCapabilityInt64:
2431       case SpvCapabilityInt64Atomics:
2432       case SpvCapabilityAtomicStorage:
2433       case SpvCapabilityInt16:
2434       case SpvCapabilityImageGatherExtended:
2435       case SpvCapabilityStorageImageMultisample:
2436       case SpvCapabilityImageCubeArray:
2437       case SpvCapabilitySampleRateShading:
2438       case SpvCapabilityInt8:
2439       case SpvCapabilityInputAttachment:
2440       case SpvCapabilitySparseResidency:
2441       case SpvCapabilityMinLod:
2442       case SpvCapabilityImageMSArray:
2443       case SpvCapabilityStorageImageExtendedFormats:
2444       case SpvCapabilityTransformFeedback:
2445       case SpvCapabilityStorageImageReadWithoutFormat:
2446       case SpvCapabilityStorageImageWriteWithoutFormat:
2447          vtn_warn("Unsupported SPIR-V capability: %s",
2448                   spirv_capability_to_string(cap));
2449          break;
2450
2451       case SpvCapabilityAddresses:
2452       case SpvCapabilityKernel:
2453       case SpvCapabilityImageBasic:
2454       case SpvCapabilityImageReadWrite:
2455       case SpvCapabilityImageMipmap:
2456       case SpvCapabilityPipes:
2457       case SpvCapabilityGroups:
2458       case SpvCapabilityDeviceEnqueue:
2459       case SpvCapabilityLiteralSampler:
2460       case SpvCapabilityGenericPointer:
2461          vtn_warn("Unsupported OpenCL-style SPIR-V capability: %s",
2462                   spirv_capability_to_string(cap));
2463          break;
2464       }
2465       break;
2466    }
2467
2468    case SpvOpExtInstImport:
2469       vtn_handle_extension(b, opcode, w, count);
2470       break;
2471
2472    case SpvOpMemoryModel:
2473       assert(w[1] == SpvAddressingModelLogical);
2474       assert(w[2] == SpvMemoryModelGLSL450);
2475       break;
2476
2477    case SpvOpEntryPoint: {
2478       struct vtn_value *entry_point = &b->values[w[2]];
2479       /* Let this be a name label regardless */
2480       unsigned name_words;
2481       entry_point->name = vtn_string_literal(b, &w[3], count - 3, &name_words);
2482
2483       if (strcmp(entry_point->name, b->entry_point_name) != 0 ||
2484           stage_for_execution_model(w[1]) != b->entry_point_stage)
2485          break;
2486
2487       assert(b->entry_point == NULL);
2488       b->entry_point = entry_point;
2489       break;
2490    }
2491
2492    case SpvOpString:
2493       vtn_push_value(b, w[1], vtn_value_type_string)->str =
2494          vtn_string_literal(b, &w[2], count - 2, NULL);
2495       break;
2496
2497    case SpvOpName:
2498       b->values[w[1]].name = vtn_string_literal(b, &w[2], count - 2, NULL);
2499       break;
2500
2501    case SpvOpMemberName:
2502       /* TODO */
2503       break;
2504
2505    case SpvOpExecutionMode:
2506    case SpvOpDecorationGroup:
2507    case SpvOpDecorate:
2508    case SpvOpMemberDecorate:
2509    case SpvOpGroupDecorate:
2510    case SpvOpGroupMemberDecorate:
2511       vtn_handle_decoration(b, opcode, w, count);
2512       break;
2513
2514    default:
2515       return false; /* End of preamble */
2516    }
2517
2518    return true;
2519 }
2520
2521 static void
2522 vtn_handle_execution_mode(struct vtn_builder *b, struct vtn_value *entry_point,
2523                           const struct vtn_decoration *mode, void *data)
2524 {
2525    assert(b->entry_point == entry_point);
2526
2527    switch(mode->exec_mode) {
2528    case SpvExecutionModeOriginUpperLeft:
2529    case SpvExecutionModeOriginLowerLeft:
2530       b->origin_upper_left =
2531          (mode->exec_mode == SpvExecutionModeOriginUpperLeft);
2532       break;
2533
2534    case SpvExecutionModeEarlyFragmentTests:
2535       assert(b->shader->stage == MESA_SHADER_FRAGMENT);
2536       b->shader->info.fs.early_fragment_tests = true;
2537       break;
2538
2539    case SpvExecutionModeInvocations:
2540       assert(b->shader->stage == MESA_SHADER_GEOMETRY);
2541       b->shader->info.gs.invocations = MAX2(1, mode->literals[0]);
2542       break;
2543
2544    case SpvExecutionModeDepthReplacing:
2545       assert(b->shader->stage == MESA_SHADER_FRAGMENT);
2546       b->shader->info.fs.depth_layout = FRAG_DEPTH_LAYOUT_ANY;
2547       break;
2548    case SpvExecutionModeDepthGreater:
2549       assert(b->shader->stage == MESA_SHADER_FRAGMENT);
2550       b->shader->info.fs.depth_layout = FRAG_DEPTH_LAYOUT_GREATER;
2551       break;
2552    case SpvExecutionModeDepthLess:
2553       assert(b->shader->stage == MESA_SHADER_FRAGMENT);
2554       b->shader->info.fs.depth_layout = FRAG_DEPTH_LAYOUT_LESS;
2555       break;
2556    case SpvExecutionModeDepthUnchanged:
2557       assert(b->shader->stage == MESA_SHADER_FRAGMENT);
2558       b->shader->info.fs.depth_layout = FRAG_DEPTH_LAYOUT_UNCHANGED;
2559       break;
2560
2561    case SpvExecutionModeLocalSize:
2562       assert(b->shader->stage == MESA_SHADER_COMPUTE);
2563       b->shader->info.cs.local_size[0] = mode->literals[0];
2564       b->shader->info.cs.local_size[1] = mode->literals[1];
2565       b->shader->info.cs.local_size[2] = mode->literals[2];
2566       break;
2567    case SpvExecutionModeLocalSizeHint:
2568       break; /* Nothing to do with this */
2569
2570    case SpvExecutionModeOutputVertices:
2571       assert(b->shader->stage == MESA_SHADER_GEOMETRY);
2572       b->shader->info.gs.vertices_out = mode->literals[0];
2573       break;
2574
2575    case SpvExecutionModeInputPoints:
2576    case SpvExecutionModeInputLines:
2577    case SpvExecutionModeInputLinesAdjacency:
2578    case SpvExecutionModeTriangles:
2579    case SpvExecutionModeInputTrianglesAdjacency:
2580    case SpvExecutionModeQuads:
2581    case SpvExecutionModeIsolines:
2582       if (b->shader->stage == MESA_SHADER_GEOMETRY) {
2583          b->shader->info.gs.vertices_in =
2584             vertices_in_from_spv_execution_mode(mode->exec_mode);
2585       } else {
2586          assert(!"Tesselation shaders not yet supported");
2587       }
2588       break;
2589
2590    case SpvExecutionModeOutputPoints:
2591    case SpvExecutionModeOutputLineStrip:
2592    case SpvExecutionModeOutputTriangleStrip:
2593       assert(b->shader->stage == MESA_SHADER_GEOMETRY);
2594       b->shader->info.gs.output_primitive =
2595          gl_primitive_from_spv_execution_mode(mode->exec_mode);
2596       break;
2597
2598    case SpvExecutionModeSpacingEqual:
2599    case SpvExecutionModeSpacingFractionalEven:
2600    case SpvExecutionModeSpacingFractionalOdd:
2601    case SpvExecutionModeVertexOrderCw:
2602    case SpvExecutionModeVertexOrderCcw:
2603    case SpvExecutionModePointMode:
2604       assert(!"TODO: Add tessellation metadata");
2605       break;
2606
2607    case SpvExecutionModePixelCenterInteger:
2608       b->pixel_center_integer = true;
2609       break;
2610
2611    case SpvExecutionModeXfb:
2612       assert(!"Unhandled execution mode");
2613       break;
2614
2615    case SpvExecutionModeVecTypeHint:
2616    case SpvExecutionModeContractionOff:
2617       break; /* OpenCL */
2618    }
2619 }
2620
2621 static bool
2622 vtn_handle_variable_or_type_instruction(struct vtn_builder *b, SpvOp opcode,
2623                                         const uint32_t *w, unsigned count)
2624 {
2625    switch (opcode) {
2626    case SpvOpSource:
2627    case SpvOpSourceContinued:
2628    case SpvOpSourceExtension:
2629    case SpvOpExtension:
2630    case SpvOpCapability:
2631    case SpvOpExtInstImport:
2632    case SpvOpMemoryModel:
2633    case SpvOpEntryPoint:
2634    case SpvOpExecutionMode:
2635    case SpvOpString:
2636    case SpvOpName:
2637    case SpvOpMemberName:
2638    case SpvOpDecorationGroup:
2639    case SpvOpDecorate:
2640    case SpvOpMemberDecorate:
2641    case SpvOpGroupDecorate:
2642    case SpvOpGroupMemberDecorate:
2643       assert(!"Invalid opcode types and variables section");
2644       break;
2645
2646    case SpvOpTypeVoid:
2647    case SpvOpTypeBool:
2648    case SpvOpTypeInt:
2649    case SpvOpTypeFloat:
2650    case SpvOpTypeVector:
2651    case SpvOpTypeMatrix:
2652    case SpvOpTypeImage:
2653    case SpvOpTypeSampler:
2654    case SpvOpTypeSampledImage:
2655    case SpvOpTypeArray:
2656    case SpvOpTypeRuntimeArray:
2657    case SpvOpTypeStruct:
2658    case SpvOpTypeOpaque:
2659    case SpvOpTypePointer:
2660    case SpvOpTypeFunction:
2661    case SpvOpTypeEvent:
2662    case SpvOpTypeDeviceEvent:
2663    case SpvOpTypeReserveId:
2664    case SpvOpTypeQueue:
2665    case SpvOpTypePipe:
2666       vtn_handle_type(b, opcode, w, count);
2667       break;
2668
2669    case SpvOpConstantTrue:
2670    case SpvOpConstantFalse:
2671    case SpvOpConstant:
2672    case SpvOpConstantComposite:
2673    case SpvOpConstantSampler:
2674    case SpvOpConstantNull:
2675    case SpvOpSpecConstantTrue:
2676    case SpvOpSpecConstantFalse:
2677    case SpvOpSpecConstant:
2678    case SpvOpSpecConstantComposite:
2679    case SpvOpSpecConstantOp:
2680       vtn_handle_constant(b, opcode, w, count);
2681       break;
2682
2683    case SpvOpVariable:
2684       vtn_handle_variables(b, opcode, w, count);
2685       break;
2686
2687    default:
2688       return false; /* End of preamble */
2689    }
2690
2691    return true;
2692 }
2693
2694 static bool
2695 vtn_handle_body_instruction(struct vtn_builder *b, SpvOp opcode,
2696                             const uint32_t *w, unsigned count)
2697 {
2698    switch (opcode) {
2699    case SpvOpLabel:
2700       break;
2701
2702    case SpvOpLoopMerge:
2703    case SpvOpSelectionMerge:
2704       /* This is handled by cfg pre-pass and walk_blocks */
2705       break;
2706
2707    case SpvOpUndef: {
2708       struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_undef);
2709       val->type = vtn_value(b, w[1], vtn_value_type_type)->type;
2710       break;
2711    }
2712
2713    case SpvOpExtInst:
2714       vtn_handle_extension(b, opcode, w, count);
2715       break;
2716
2717    case SpvOpVariable:
2718    case SpvOpLoad:
2719    case SpvOpStore:
2720    case SpvOpCopyMemory:
2721    case SpvOpCopyMemorySized:
2722    case SpvOpAccessChain:
2723    case SpvOpInBoundsAccessChain:
2724    case SpvOpArrayLength:
2725       vtn_handle_variables(b, opcode, w, count);
2726       break;
2727
2728    case SpvOpFunctionCall:
2729       vtn_handle_function_call(b, opcode, w, count);
2730       break;
2731
2732    case SpvOpSampledImage:
2733    case SpvOpImage:
2734    case SpvOpImageSampleImplicitLod:
2735    case SpvOpImageSampleExplicitLod:
2736    case SpvOpImageSampleDrefImplicitLod:
2737    case SpvOpImageSampleDrefExplicitLod:
2738    case SpvOpImageSampleProjImplicitLod:
2739    case SpvOpImageSampleProjExplicitLod:
2740    case SpvOpImageSampleProjDrefImplicitLod:
2741    case SpvOpImageSampleProjDrefExplicitLod:
2742    case SpvOpImageFetch:
2743    case SpvOpImageGather:
2744    case SpvOpImageDrefGather:
2745    case SpvOpImageQuerySizeLod:
2746    case SpvOpImageQueryLod:
2747    case SpvOpImageQueryLevels:
2748    case SpvOpImageQuerySamples:
2749       vtn_handle_texture(b, opcode, w, count);
2750       break;
2751
2752    case SpvOpImageRead:
2753    case SpvOpImageWrite:
2754    case SpvOpImageTexelPointer:
2755       vtn_handle_image(b, opcode, w, count);
2756       break;
2757
2758    case SpvOpImageQuerySize: {
2759       struct vtn_access_chain *image =
2760          vtn_value(b, w[3], vtn_value_type_access_chain)->access_chain;
2761       if (glsl_type_is_image(image->var->var->interface_type)) {
2762          vtn_handle_image(b, opcode, w, count);
2763       } else {
2764          vtn_handle_texture(b, opcode, w, count);
2765       }
2766       break;
2767    }
2768
2769    case SpvOpAtomicLoad:
2770    case SpvOpAtomicExchange:
2771    case SpvOpAtomicCompareExchange:
2772    case SpvOpAtomicCompareExchangeWeak:
2773    case SpvOpAtomicIIncrement:
2774    case SpvOpAtomicIDecrement:
2775    case SpvOpAtomicIAdd:
2776    case SpvOpAtomicISub:
2777    case SpvOpAtomicSMin:
2778    case SpvOpAtomicUMin:
2779    case SpvOpAtomicSMax:
2780    case SpvOpAtomicUMax:
2781    case SpvOpAtomicAnd:
2782    case SpvOpAtomicOr:
2783    case SpvOpAtomicXor: {
2784       struct vtn_value *pointer = vtn_untyped_value(b, w[3]);
2785       if (pointer->value_type == vtn_value_type_image_pointer) {
2786          vtn_handle_image(b, opcode, w, count);
2787       } else {
2788          assert(pointer->value_type == vtn_value_type_access_chain);
2789          vtn_handle_ssbo_or_shared_atomic(b, opcode, w, count);
2790       }
2791       break;
2792    }
2793
2794    case SpvOpAtomicStore: {
2795       struct vtn_value *pointer = vtn_untyped_value(b, w[1]);
2796       if (pointer->value_type == vtn_value_type_image_pointer) {
2797          vtn_handle_image(b, opcode, w, count);
2798       } else {
2799          assert(pointer->value_type == vtn_value_type_access_chain);
2800          vtn_handle_ssbo_or_shared_atomic(b, opcode, w, count);
2801       }
2802       break;
2803    }
2804
2805    case SpvOpSNegate:
2806    case SpvOpFNegate:
2807    case SpvOpNot:
2808    case SpvOpAny:
2809    case SpvOpAll:
2810    case SpvOpConvertFToU:
2811    case SpvOpConvertFToS:
2812    case SpvOpConvertSToF:
2813    case SpvOpConvertUToF:
2814    case SpvOpUConvert:
2815    case SpvOpSConvert:
2816    case SpvOpFConvert:
2817    case SpvOpQuantizeToF16:
2818    case SpvOpConvertPtrToU:
2819    case SpvOpConvertUToPtr:
2820    case SpvOpPtrCastToGeneric:
2821    case SpvOpGenericCastToPtr:
2822    case SpvOpBitcast:
2823    case SpvOpIsNan:
2824    case SpvOpIsInf:
2825    case SpvOpIsFinite:
2826    case SpvOpIsNormal:
2827    case SpvOpSignBitSet:
2828    case SpvOpLessOrGreater:
2829    case SpvOpOrdered:
2830    case SpvOpUnordered:
2831    case SpvOpIAdd:
2832    case SpvOpFAdd:
2833    case SpvOpISub:
2834    case SpvOpFSub:
2835    case SpvOpIMul:
2836    case SpvOpFMul:
2837    case SpvOpUDiv:
2838    case SpvOpSDiv:
2839    case SpvOpFDiv:
2840    case SpvOpUMod:
2841    case SpvOpSRem:
2842    case SpvOpSMod:
2843    case SpvOpFRem:
2844    case SpvOpFMod:
2845    case SpvOpVectorTimesScalar:
2846    case SpvOpDot:
2847    case SpvOpIAddCarry:
2848    case SpvOpISubBorrow:
2849    case SpvOpUMulExtended:
2850    case SpvOpSMulExtended:
2851    case SpvOpShiftRightLogical:
2852    case SpvOpShiftRightArithmetic:
2853    case SpvOpShiftLeftLogical:
2854    case SpvOpLogicalEqual:
2855    case SpvOpLogicalNotEqual:
2856    case SpvOpLogicalOr:
2857    case SpvOpLogicalAnd:
2858    case SpvOpLogicalNot:
2859    case SpvOpBitwiseOr:
2860    case SpvOpBitwiseXor:
2861    case SpvOpBitwiseAnd:
2862    case SpvOpSelect:
2863    case SpvOpIEqual:
2864    case SpvOpFOrdEqual:
2865    case SpvOpFUnordEqual:
2866    case SpvOpINotEqual:
2867    case SpvOpFOrdNotEqual:
2868    case SpvOpFUnordNotEqual:
2869    case SpvOpULessThan:
2870    case SpvOpSLessThan:
2871    case SpvOpFOrdLessThan:
2872    case SpvOpFUnordLessThan:
2873    case SpvOpUGreaterThan:
2874    case SpvOpSGreaterThan:
2875    case SpvOpFOrdGreaterThan:
2876    case SpvOpFUnordGreaterThan:
2877    case SpvOpULessThanEqual:
2878    case SpvOpSLessThanEqual:
2879    case SpvOpFOrdLessThanEqual:
2880    case SpvOpFUnordLessThanEqual:
2881    case SpvOpUGreaterThanEqual:
2882    case SpvOpSGreaterThanEqual:
2883    case SpvOpFOrdGreaterThanEqual:
2884    case SpvOpFUnordGreaterThanEqual:
2885    case SpvOpDPdx:
2886    case SpvOpDPdy:
2887    case SpvOpFwidth:
2888    case SpvOpDPdxFine:
2889    case SpvOpDPdyFine:
2890    case SpvOpFwidthFine:
2891    case SpvOpDPdxCoarse:
2892    case SpvOpDPdyCoarse:
2893    case SpvOpFwidthCoarse:
2894    case SpvOpBitFieldInsert:
2895    case SpvOpBitFieldSExtract:
2896    case SpvOpBitFieldUExtract:
2897    case SpvOpBitReverse:
2898    case SpvOpBitCount:
2899    case SpvOpTranspose:
2900    case SpvOpOuterProduct:
2901    case SpvOpMatrixTimesScalar:
2902    case SpvOpVectorTimesMatrix:
2903    case SpvOpMatrixTimesVector:
2904    case SpvOpMatrixTimesMatrix:
2905       vtn_handle_alu(b, opcode, w, count);
2906       break;
2907
2908    case SpvOpVectorExtractDynamic:
2909    case SpvOpVectorInsertDynamic:
2910    case SpvOpVectorShuffle:
2911    case SpvOpCompositeConstruct:
2912    case SpvOpCompositeExtract:
2913    case SpvOpCompositeInsert:
2914    case SpvOpCopyObject:
2915       vtn_handle_composite(b, opcode, w, count);
2916       break;
2917
2918    case SpvOpEmitVertex:
2919    case SpvOpEndPrimitive:
2920    case SpvOpEmitStreamVertex:
2921    case SpvOpEndStreamPrimitive:
2922    case SpvOpControlBarrier:
2923    case SpvOpMemoryBarrier:
2924       vtn_handle_barrier(b, opcode, w, count);
2925       break;
2926
2927    default:
2928       unreachable("Unhandled opcode");
2929    }
2930
2931    return true;
2932 }
2933
2934 nir_function *
2935 spirv_to_nir(const uint32_t *words, size_t word_count,
2936              struct nir_spirv_specialization *spec, unsigned num_spec,
2937              gl_shader_stage stage, const char *entry_point_name,
2938              const nir_shader_compiler_options *options)
2939 {
2940    const uint32_t *word_end = words + word_count;
2941
2942    /* Handle the SPIR-V header (first 4 dwords)  */
2943    assert(word_count > 5);
2944
2945    assert(words[0] == SpvMagicNumber);
2946    assert(words[1] >= 0x10000);
2947    /* words[2] == generator magic */
2948    unsigned value_id_bound = words[3];
2949    assert(words[4] == 0);
2950
2951    words+= 5;
2952
2953    /* Initialize the stn_builder object */
2954    struct vtn_builder *b = rzalloc(NULL, struct vtn_builder);
2955    b->value_id_bound = value_id_bound;
2956    b->values = rzalloc_array(b, struct vtn_value, value_id_bound);
2957    exec_list_make_empty(&b->functions);
2958    b->entry_point_stage = stage;
2959    b->entry_point_name = entry_point_name;
2960
2961    /* Handle all the preamble instructions */
2962    words = vtn_foreach_instruction(b, words, word_end,
2963                                    vtn_handle_preamble_instruction);
2964
2965    if (b->entry_point == NULL) {
2966       assert(!"Entry point not found");
2967       ralloc_free(b);
2968       return NULL;
2969    }
2970
2971    b->shader = nir_shader_create(NULL, stage, options);
2972
2973    /* Set shader info defaults */
2974    b->shader->info.gs.invocations = 1;
2975
2976    /* Parse execution modes */
2977    vtn_foreach_execution_mode(b, b->entry_point,
2978                               vtn_handle_execution_mode, NULL);
2979
2980    b->specializations = spec;
2981    b->num_specializations = num_spec;
2982
2983    /* Handle all variable, type, and constant instructions */
2984    words = vtn_foreach_instruction(b, words, word_end,
2985                                    vtn_handle_variable_or_type_instruction);
2986
2987    vtn_build_cfg(b, words, word_end);
2988
2989    foreach_list_typed(struct vtn_function, func, node, &b->functions) {
2990       b->impl = func->impl;
2991       b->const_table = _mesa_hash_table_create(b, _mesa_hash_pointer,
2992                                                _mesa_key_pointer_equal);
2993
2994       vtn_function_emit(b, func, vtn_handle_body_instruction);
2995    }
2996
2997    assert(b->entry_point->value_type == vtn_value_type_function);
2998    nir_function *entry_point = b->entry_point->func->impl->function;
2999    assert(entry_point);
3000
3001    ralloc_free(b);
3002
3003    return entry_point;
3004 }