[mesa.git] / src / compiler / spirv / spirv_to_nir.c

/*
 * Copyright © 2015 Intel Corporation
 *
 * Permission is hereby granted, free of charge, to any person obtaining a
 * copy of this software and associated documentation files (the "Software"),
 * to deal in the Software without restriction, including without limitation
 * the rights to use, copy, modify, merge, publish, distribute, sublicense,
 * and/or sell copies of the Software, and to permit persons to whom the
 * Software is furnished to do so, subject to the following conditions:
 *
 * The above copyright notice and this permission notice (including the next
 * paragraph) shall be included in all copies or substantial portions of the
 * Software.
 *
 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL
 * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
 * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
 * IN THE SOFTWARE.
 *
 * Authors:
 *    Jason Ekstrand (jason@jlekstrand.net)
 *
 */

#include "vtn_private.h"
#include "nir/nir_vla.h"
#include "nir/nir_control_flow.h"
#include "nir/nir_constant_expressions.h"
#include "nir/nir_deref.h"
#include "spirv_info.h"

#include "util/format/u_format.h"
#include "util/u_math.h"

#include <stdio.h>

void
vtn_log(struct vtn_builder *b, enum nir_spirv_debug_level level,
        size_t spirv_offset, const char *message)
{
   if (b->options->debug.func) {
      b->options->debug.func(b->options->debug.private_data,
                             level, spirv_offset, message);
   }

#ifndef NDEBUG
   if (level >= NIR_SPIRV_DEBUG_LEVEL_WARNING)
      fprintf(stderr, "%s\n", message);
#endif
}

void
vtn_logf(struct vtn_builder *b, enum nir_spirv_debug_level level,
         size_t spirv_offset, const char *fmt, ...)
{
   va_list args;
   char *msg;

   va_start(args, fmt);
   msg = ralloc_vasprintf(NULL, fmt, args);
   va_end(args);

   vtn_log(b, level, spirv_offset, msg);

   ralloc_free(msg);
}

static void
vtn_log_err(struct vtn_builder *b,
            enum nir_spirv_debug_level level, const char *prefix,
            const char *file, unsigned line,
            const char *fmt, va_list args)
{
   char *msg;

   msg = ralloc_strdup(NULL, prefix);

#ifndef NDEBUG
   ralloc_asprintf_append(&msg, "    In file %s:%u\n", file, line);
#endif

   ralloc_asprintf_append(&msg, "    ");

   ralloc_vasprintf_append(&msg, fmt, args);

   ralloc_asprintf_append(&msg, "\n    %zu bytes into the SPIR-V binary",
                          b->spirv_offset);

   if (b->file) {
      ralloc_asprintf_append(&msg,
                             "\n    in SPIR-V source file %s, line %d, col %d",
                             b->file, b->line, b->col);
   }

   vtn_log(b, level, b->spirv_offset, msg);

   ralloc_free(msg);
}

static void
vtn_dump_shader(struct vtn_builder *b, const char *path, const char *prefix)
{
   static int idx = 0;

   char filename[1024];
   int len = snprintf(filename, sizeof(filename), "%s/%s-%d.spirv",
                      path, prefix, idx++);
   if (len < 0 || len >= sizeof(filename))
      return;

   FILE *f = fopen(filename, "w");
   if (f == NULL)
      return;

   fwrite(b->spirv, sizeof(*b->spirv), b->spirv_word_count, f);
   fclose(f);

   vtn_info("SPIR-V shader dumped to %s", filename);
}

void
_vtn_warn(struct vtn_builder *b, const char *file, unsigned line,
          const char *fmt, ...)
{
   va_list args;

   va_start(args, fmt);
   vtn_log_err(b, NIR_SPIRV_DEBUG_LEVEL_WARNING, "SPIR-V WARNING:\n",
               file, line, fmt, args);
   va_end(args);
}

void
_vtn_err(struct vtn_builder *b, const char *file, unsigned line,
          const char *fmt, ...)
{
   va_list args;

   va_start(args, fmt);
   vtn_log_err(b, NIR_SPIRV_DEBUG_LEVEL_ERROR, "SPIR-V ERROR:\n",
               file, line, fmt, args);
   va_end(args);
}

void
_vtn_fail(struct vtn_builder *b, const char *file, unsigned line,
          const char *fmt, ...)
{
   va_list args;

   va_start(args, fmt);
   vtn_log_err(b, NIR_SPIRV_DEBUG_LEVEL_ERROR, "SPIR-V parsing FAILED:\n",
               file, line, fmt, args);
   va_end(args);

   const char *dump_path = getenv("MESA_SPIRV_FAIL_DUMP_PATH");
   if (dump_path)
      vtn_dump_shader(b, dump_path, "fail");

   longjmp(b->fail_jump, 1);
}

static struct vtn_ssa_value *
vtn_undef_ssa_value(struct vtn_builder *b, const struct glsl_type *type)
{
   struct vtn_ssa_value *val = rzalloc(b, struct vtn_ssa_value);
   val->type = glsl_get_bare_type(type);

   if (glsl_type_is_vector_or_scalar(type)) {
      unsigned num_components = glsl_get_vector_elements(val->type);
      unsigned bit_size = glsl_get_bit_size(val->type);
      val->def = nir_ssa_undef(&b->nb, num_components, bit_size);
   } else {
      unsigned elems = glsl_get_length(val->type);
      val->elems = ralloc_array(b, struct vtn_ssa_value *, elems);
      if (glsl_type_is_array_or_matrix(type)) {
         const struct glsl_type *elem_type = glsl_get_array_element(type);
         for (unsigned i = 0; i < elems; i++)
            val->elems[i] = vtn_undef_ssa_value(b, elem_type);
      } else {
         vtn_assert(glsl_type_is_struct_or_ifc(type));
         for (unsigned i = 0; i < elems; i++) {
            const struct glsl_type *elem_type = glsl_get_struct_field(type, i);
            val->elems[i] = vtn_undef_ssa_value(b, elem_type);
         }
      }
   }

   return val;
}

static struct vtn_ssa_value *
vtn_const_ssa_value(struct vtn_builder *b, nir_constant *constant,
                    const struct glsl_type *type)
{
   struct hash_entry *entry = _mesa_hash_table_search(b->const_table, constant);

   if (entry)
      return entry->data;

   struct vtn_ssa_value *val = rzalloc(b, struct vtn_ssa_value);
   val->type = glsl_get_bare_type(type);

   if (glsl_type_is_vector_or_scalar(type)) {
      unsigned num_components = glsl_get_vector_elements(val->type);
      unsigned bit_size = glsl_get_bit_size(type);
      nir_load_const_instr *load =
         nir_load_const_instr_create(b->shader, num_components, bit_size);

      memcpy(load->value, constant->values,
             sizeof(nir_const_value) * num_components);

      nir_instr_insert_before_cf_list(&b->nb.impl->body, &load->instr);
      val->def = &load->def;
   } else {
      unsigned elems = glsl_get_length(val->type);
      val->elems = ralloc_array(b, struct vtn_ssa_value *, elems);
      if (glsl_type_is_array_or_matrix(type)) {
         const struct glsl_type *elem_type = glsl_get_array_element(type);
         for (unsigned i = 0; i < elems; i++) {
            val->elems[i] = vtn_const_ssa_value(b, constant->elements[i],
                                                elem_type);
         }
      } else {
         vtn_assert(glsl_type_is_struct_or_ifc(type));
         for (unsigned i = 0; i < elems; i++) {
            const struct glsl_type *elem_type = glsl_get_struct_field(type, i);
            val->elems[i] = vtn_const_ssa_value(b, constant->elements[i],
                                                elem_type);
         }
      }
   }

   return val;
}

struct vtn_ssa_value *
vtn_ssa_value(struct vtn_builder *b, uint32_t value_id)
{
   struct vtn_value *val = vtn_untyped_value(b, value_id);
   switch (val->value_type) {
   case vtn_value_type_undef:
      return vtn_undef_ssa_value(b, val->type->type);

   case vtn_value_type_constant:
      return vtn_const_ssa_value(b, val->constant, val->type->type);

   case vtn_value_type_ssa:
      return val->ssa;

   case vtn_value_type_pointer:
      vtn_assert(val->pointer->ptr_type && val->pointer->ptr_type->type);
      struct vtn_ssa_value *ssa =
         vtn_create_ssa_value(b, val->pointer->ptr_type->type);
      ssa->def = vtn_pointer_to_ssa(b, val->pointer);
      return ssa;

   default:
      vtn_fail("Invalid type for an SSA value");
   }
}

struct vtn_value *
vtn_push_ssa_value(struct vtn_builder *b, uint32_t value_id,
                   struct vtn_ssa_value *ssa)
{
   struct vtn_type *type = vtn_get_value_type(b, value_id);

   /* See vtn_create_ssa_value */
   vtn_fail_if(ssa->type != glsl_get_bare_type(type->type),
               "Type mismatch for SPIR-V SSA value");

   struct vtn_value *val;
   if (type->base_type == vtn_base_type_pointer) {
      val = vtn_push_pointer(b, value_id, vtn_pointer_from_ssa(b, ssa->def, type));
   } else {
      /* Don't trip the value_type_ssa check in vtn_push_value */
      val = vtn_push_value(b, value_id, vtn_value_type_invalid);
      val->value_type = vtn_value_type_ssa;
      val->ssa = ssa;
   }

   return val;
}

nir_ssa_def *
vtn_get_nir_ssa(struct vtn_builder *b, uint32_t value_id)
{
   struct vtn_ssa_value *ssa = vtn_ssa_value(b, value_id);
   vtn_fail_if(!glsl_type_is_vector_or_scalar(ssa->type),
               "Expected a vector or scalar type");
   return ssa->def;
}

struct vtn_value *
vtn_push_nir_ssa(struct vtn_builder *b, uint32_t value_id, nir_ssa_def *def)
{
   /* Types for all SPIR-V SSA values are set as part of a pre-pass so the
    * type will be valid by the time we get here.
    */
   struct vtn_type *type = vtn_get_value_type(b, value_id);
   vtn_fail_if(def->num_components != glsl_get_vector_elements(type->type) ||
               def->bit_size != glsl_get_bit_size(type->type),
               "Mismatch between NIR and SPIR-V type.");
   struct vtn_ssa_value *ssa = vtn_create_ssa_value(b, type->type);
   ssa->def = def;
   return vtn_push_ssa_value(b, value_id, ssa);
}

static nir_deref_instr *
vtn_get_image(struct vtn_builder *b, uint32_t value_id)
{
   struct vtn_type *type = vtn_get_value_type(b, value_id);
   vtn_assert(type->base_type == vtn_base_type_image);
   return nir_build_deref_cast(&b->nb, vtn_get_nir_ssa(b, value_id),
                               nir_var_uniform, type->glsl_image, 0);
}

static void
vtn_push_image(struct vtn_builder *b, uint32_t value_id,
               nir_deref_instr *deref)
{
   struct vtn_type *type = vtn_get_value_type(b, value_id);
   vtn_assert(type->base_type == vtn_base_type_image);
   vtn_push_nir_ssa(b, value_id, &deref->dest.ssa);
}

static nir_deref_instr *
vtn_get_sampler(struct vtn_builder *b, uint32_t value_id)
{
   struct vtn_type *type = vtn_get_value_type(b, value_id);
   vtn_assert(type->base_type == vtn_base_type_sampler);
   return nir_build_deref_cast(&b->nb, vtn_get_nir_ssa(b, value_id),
                               nir_var_uniform, glsl_bare_sampler_type(), 0);
}

nir_ssa_def *
vtn_sampled_image_to_nir_ssa(struct vtn_builder *b,
                             struct vtn_sampled_image si)
{
   return nir_vec2(&b->nb, &si.image->dest.ssa, &si.sampler->dest.ssa);
}

static void
vtn_push_sampled_image(struct vtn_builder *b, uint32_t value_id,
                       struct vtn_sampled_image si)
{
   struct vtn_type *type = vtn_get_value_type(b, value_id);
   vtn_assert(type->base_type == vtn_base_type_sampled_image);
   vtn_push_nir_ssa(b, value_id, vtn_sampled_image_to_nir_ssa(b, si));
}

static struct vtn_sampled_image
vtn_get_sampled_image(struct vtn_builder *b, uint32_t value_id)
{
   struct vtn_type *type = vtn_get_value_type(b, value_id);
   vtn_assert(type->base_type == vtn_base_type_sampled_image);
   nir_ssa_def *si_vec2 = vtn_get_nir_ssa(b, value_id);

   struct vtn_sampled_image si = { NULL, };
   si.image = nir_build_deref_cast(&b->nb, nir_channel(&b->nb, si_vec2, 0),
                                   nir_var_uniform,
                                   type->image->glsl_image, 0);
   si.sampler = nir_build_deref_cast(&b->nb, nir_channel(&b->nb, si_vec2, 1),
                                     nir_var_uniform,
                                     glsl_bare_sampler_type(), 0);
   return si;
}

static char *
vtn_string_literal(struct vtn_builder *b, const uint32_t *words,
                   unsigned word_count, unsigned *words_used)
{
   char *dup = ralloc_strndup(b, (char *)words, word_count * sizeof(*words));
   if (words_used) {
      /* Ammount of space taken by the string (including the null) */
      unsigned len = strlen(dup) + 1;
      *words_used = DIV_ROUND_UP(len, sizeof(*words));
   }
   return dup;
}

const uint32_t *
vtn_foreach_instruction(struct vtn_builder *b, const uint32_t *start,
                        const uint32_t *end, vtn_instruction_handler handler)
{
   b->file = NULL;
   b->line = -1;
   b->col = -1;

   const uint32_t *w = start;
   while (w < end) {
      SpvOp opcode = w[0] & SpvOpCodeMask;
      unsigned count = w[0] >> SpvWordCountShift;
      vtn_assert(count >= 1 && w + count <= end);

      b->spirv_offset = (uint8_t *)w - (uint8_t *)b->spirv;

      switch (opcode) {
      case SpvOpNop:
         break; /* Do nothing */

      case SpvOpLine:
         b->file = vtn_value(b, w[1], vtn_value_type_string)->str;
         b->line = w[2];
         b->col = w[3];
         break;

      case SpvOpNoLine:
         b->file = NULL;
         b->line = -1;
         b->col = -1;
         break;

      default:
         if (!handler(b, opcode, w, count))
            return w;
         break;
      }

      w += count;
   }

   b->spirv_offset = 0;
   b->file = NULL;
   b->line = -1;
   b->col = -1;

   assert(w == end);
   return w;
}

static bool
vtn_handle_non_semantic_instruction(struct vtn_builder *b, SpvOp ext_opcode,
                                    const uint32_t *w, unsigned count)
{
   /* Do nothing. */
   return true;
}

static void
vtn_handle_extension(struct vtn_builder *b, SpvOp opcode,
                     const uint32_t *w, unsigned count)
{
   const char *ext = (const char *)&w[2];
   switch (opcode) {
   case SpvOpExtInstImport: {
      struct vtn_value *val = vtn_push_value(b, w[1], vtn_value_type_extension);
      if (strcmp(ext, "GLSL.std.450") == 0) {
         val->ext_handler = vtn_handle_glsl450_instruction;
      } else if ((strcmp(ext, "SPV_AMD_gcn_shader") == 0)
                && (b->options && b->options->caps.amd_gcn_shader)) {
         val->ext_handler = vtn_handle_amd_gcn_shader_instruction;
      } else if ((strcmp(ext, "SPV_AMD_shader_ballot") == 0)
                && (b->options && b->options->caps.amd_shader_ballot)) {
         val->ext_handler = vtn_handle_amd_shader_ballot_instruction;
      } else if ((strcmp(ext, "SPV_AMD_shader_trinary_minmax") == 0)
                && (b->options && b->options->caps.amd_trinary_minmax)) {
         val->ext_handler = vtn_handle_amd_shader_trinary_minmax_instruction;
      } else if ((strcmp(ext, "SPV_AMD_shader_explicit_vertex_parameter") == 0)
                && (b->options && b->options->caps.amd_shader_explicit_vertex_parameter)) {
         val->ext_handler = vtn_handle_amd_shader_explicit_vertex_parameter_instruction;
      } else if (strcmp(ext, "OpenCL.std") == 0) {
         val->ext_handler = vtn_handle_opencl_instruction;
      } else if (strstr(ext, "NonSemantic.") == ext) {
         val->ext_handler = vtn_handle_non_semantic_instruction;
      } else {
         vtn_fail("Unsupported extension: %s", ext);
      }
      break;
   }

   case SpvOpExtInst: {
      struct vtn_value *val = vtn_value(b, w[3], vtn_value_type_extension);
      bool handled = val->ext_handler(b, w[4], w, count);
      vtn_assert(handled);
      break;
   }

   default:
      vtn_fail_with_opcode("Unhandled opcode", opcode);
   }
}

static void
_foreach_decoration_helper(struct vtn_builder *b,
                           struct vtn_value *base_value,
                           int parent_member,
                           struct vtn_value *value,
                           vtn_decoration_foreach_cb cb, void *data)
{
   for (struct vtn_decoration *dec = value->decoration; dec; dec = dec->next) {
      int member;
      if (dec->scope == VTN_DEC_DECORATION) {
         member = parent_member;
      } else if (dec->scope >= VTN_DEC_STRUCT_MEMBER0) {
         vtn_fail_if(value->value_type != vtn_value_type_type ||
                     value->type->base_type != vtn_base_type_struct,
                     "OpMemberDecorate and OpGroupMemberDecorate are only "
                     "allowed on OpTypeStruct");
         /* This means we haven't recursed yet */
         assert(value == base_value);

         member = dec->scope - VTN_DEC_STRUCT_MEMBER0;

         vtn_fail_if(member >= base_value->type->length,
                     "OpMemberDecorate specifies member %d but the "
                     "OpTypeStruct has only %u members",
                     member, base_value->type->length);
      } else {
         /* Not a decoration */
         assert(dec->scope == VTN_DEC_EXECUTION_MODE);
         continue;
      }

      if (dec->group) {
         assert(dec->group->value_type == vtn_value_type_decoration_group);
         _foreach_decoration_helper(b, base_value, member, dec->group,
                                    cb, data);
      } else {
         cb(b, base_value, member, dec, data);
      }
   }
}

/** Iterates (recursively if needed) over all of the decorations on a value
 *
 * This function iterates over all of the decorations applied to a given
 * value.  If it encounters a decoration group, it recurses into the group
 * and iterates over all of those decorations as well.
 */
void
vtn_foreach_decoration(struct vtn_builder *b, struct vtn_value *value,
                       vtn_decoration_foreach_cb cb, void *data)
{
   _foreach_decoration_helper(b, value, -1, value, cb, data);
}

void
vtn_foreach_execution_mode(struct vtn_builder *b, struct vtn_value *value,
                           vtn_execution_mode_foreach_cb cb, void *data)
{
   for (struct vtn_decoration *dec = value->decoration; dec; dec = dec->next) {
      if (dec->scope != VTN_DEC_EXECUTION_MODE)
         continue;

      assert(dec->group == NULL);
      cb(b, value, dec, data);
   }
}

void
vtn_handle_decoration(struct vtn_builder *b, SpvOp opcode,
                      const uint32_t *w, unsigned count)
{
   const uint32_t *w_end = w + count;
   const uint32_t target = w[1];
   w += 2;

   switch (opcode) {
   case SpvOpDecorationGroup:
      vtn_push_value(b, target, vtn_value_type_decoration_group);
      break;

   case SpvOpDecorate:
   case SpvOpDecorateId:
   case SpvOpMemberDecorate:
   case SpvOpDecorateString:
   case SpvOpMemberDecorateString:
   case SpvOpExecutionMode:
   case SpvOpExecutionModeId: {
      struct vtn_value *val = vtn_untyped_value(b, target);

      struct vtn_decoration *dec = rzalloc(b, struct vtn_decoration);
      switch (opcode) {
      case SpvOpDecorate:
      case SpvOpDecorateId:
      case SpvOpDecorateString:
         dec->scope = VTN_DEC_DECORATION;
         break;
      case SpvOpMemberDecorate:
      case SpvOpMemberDecorateString:
         dec->scope = VTN_DEC_STRUCT_MEMBER0 + *(w++);
         vtn_fail_if(dec->scope < VTN_DEC_STRUCT_MEMBER0, /* overflow */
                     "Member argument of OpMemberDecorate too large");
         break;
      case SpvOpExecutionMode:
      case SpvOpExecutionModeId:
         dec->scope = VTN_DEC_EXECUTION_MODE;
         break;
      default:
         unreachable("Invalid decoration opcode");
      }
      dec->decoration = *(w++);
      dec->operands = w;

      /* Link into the list */
      dec->next = val->decoration;
      val->decoration = dec;
      break;
   }

   case SpvOpGroupMemberDecorate:
   case SpvOpGroupDecorate: {
      struct vtn_value *group =
         vtn_value(b, target, vtn_value_type_decoration_group);

      for (; w < w_end; w++) {
         struct vtn_value *val = vtn_untyped_value(b, *w);
         struct vtn_decoration *dec = rzalloc(b, struct vtn_decoration);

         dec->group = group;
         if (opcode == SpvOpGroupDecorate) {
            dec->scope = VTN_DEC_DECORATION;
         } else {
            dec->scope = VTN_DEC_STRUCT_MEMBER0 + *(++w);
            vtn_fail_if(dec->scope < 0, /* Check for overflow */
                        "Member argument of OpGroupMemberDecorate too large");
         }

         /* Link into the list */
         dec->next = val->decoration;
         val->decoration = dec;
      }
      break;
   }

   default:
      unreachable("Unhandled opcode");
   }
}

struct member_decoration_ctx {
   unsigned num_fields;
   struct glsl_struct_field *fields;
   struct vtn_type *type;
};

/**
 * Returns true if the given type contains a struct decorated Block or
 * BufferBlock
 */
bool
vtn_type_contains_block(struct vtn_builder *b, struct vtn_type *type)
{
   switch (type->base_type) {
   case vtn_base_type_array:
      return vtn_type_contains_block(b, type->array_element);
   case vtn_base_type_struct:
      if (type->block || type->buffer_block)
         return true;
      for (unsigned i = 0; i < type->length; i++) {
         if (vtn_type_contains_block(b, type->members[i]))
            return true;
      }
      return false;
   default:
      return false;
   }
}

/** Returns true if two types are "compatible", i.e. you can do an OpLoad,
 * OpStore, or OpCopyMemory between them without breaking anything.
 * Technically, the SPIR-V rules require the exact same type ID but this lets
 * us internally be a bit looser.
 */
bool
vtn_types_compatible(struct vtn_builder *b,
                     struct vtn_type *t1, struct vtn_type *t2)
{
   if (t1->id == t2->id)
      return true;

   if (t1->base_type != t2->base_type)
      return false;

   switch (t1->base_type) {
   case vtn_base_type_void:
   case vtn_base_type_scalar:
   case vtn_base_type_vector:
   case vtn_base_type_matrix:
   case vtn_base_type_image:
   case vtn_base_type_sampler:
   case vtn_base_type_sampled_image:
      return t1->type == t2->type;

   case vtn_base_type_array:
      return t1->length == t2->length &&
             vtn_types_compatible(b, t1->array_element, t2->array_element);

   case vtn_base_type_pointer:
      return vtn_types_compatible(b, t1->deref, t2->deref);

   case vtn_base_type_struct:
      if (t1->length != t2->length)
         return false;

      for (unsigned i = 0; i < t1->length; i++) {
         if (!vtn_types_compatible(b, t1->members[i], t2->members[i]))
            return false;
      }
      return true;

   case vtn_base_type_function:
      /* This case shouldn't get hit since you can't copy around function
       * types.  Just require them to be identical.
       */
      return false;
   }

   vtn_fail("Invalid base type");
}

struct vtn_type *
vtn_type_without_array(struct vtn_type *type)
{
   while (type->base_type == vtn_base_type_array)
      type = type->array_element;
   return type;
}

/* does a shallow copy of a vtn_type */

static struct vtn_type *
vtn_type_copy(struct vtn_builder *b, struct vtn_type *src)
{
   struct vtn_type *dest = ralloc(b, struct vtn_type);
   *dest = *src;

   switch (src->base_type) {
   case vtn_base_type_void:
   case vtn_base_type_scalar:
   case vtn_base_type_vector:
   case vtn_base_type_matrix:
   case vtn_base_type_array:
   case vtn_base_type_pointer:
   case vtn_base_type_image:
   case vtn_base_type_sampler:
   case vtn_base_type_sampled_image:
      /* Nothing more to do */
      break;

   case vtn_base_type_struct:
      dest->members = ralloc_array(b, struct vtn_type *, src->length);
      memcpy(dest->members, src->members,
             src->length * sizeof(src->members[0]));

      dest->offsets = ralloc_array(b, unsigned, src->length);
      memcpy(dest->offsets, src->offsets,
             src->length * sizeof(src->offsets[0]));
      break;

   case vtn_base_type_function:
      dest->params = ralloc_array(b, struct vtn_type *, src->length);
      memcpy(dest->params, src->params, src->length * sizeof(src->params[0]));
      break;
   }

   return dest;
}

static const struct glsl_type *
wrap_type_in_array(const struct glsl_type *type,
                   const struct glsl_type *array_type)
{
   if (!glsl_type_is_array(array_type))
      return type;

   const struct glsl_type *elem_type =
      wrap_type_in_array(type, glsl_get_array_element(array_type));
   return glsl_array_type(elem_type, glsl_get_length(array_type),
                          glsl_get_explicit_stride(array_type));
}

static bool
vtn_type_needs_explicit_layout(struct vtn_builder *b, enum vtn_variable_mode mode)
{
   /* For OpenCL we never want to strip the info from the types, and it makes
    * type comparisons easier in later stages.
    */
   if (b->options->environment == NIR_SPIRV_OPENCL)
      return true;

   switch (mode) {
   case vtn_variable_mode_input:
   case vtn_variable_mode_output:
      /* Layout decorations kept because we need offsets for XFB arrays of
       * blocks.
       */
      return b->shader->info.has_transform_feedback_varyings;

   case vtn_variable_mode_ssbo:
   case vtn_variable_mode_phys_ssbo:
   case vtn_variable_mode_ubo:
      return true;

   default:
      return false;
   }
}

const struct glsl_type *
vtn_type_get_nir_type(struct vtn_builder *b, struct vtn_type *type,
                      enum vtn_variable_mode mode)
{
   if (mode == vtn_variable_mode_atomic_counter) {
      vtn_fail_if(glsl_without_array(type->type) != glsl_uint_type(),
                  "Variables in the AtomicCounter storage class should be "
                  "(possibly arrays of arrays of) uint.");
      return wrap_type_in_array(glsl_atomic_uint_type(), type->type);
   }

   if (mode == vtn_variable_mode_uniform) {
      switch (type->base_type) {
      case vtn_base_type_array: {
         const struct glsl_type *elem_type =
            vtn_type_get_nir_type(b, type->array_element, mode);

         return glsl_array_type(elem_type, type->length,
                                glsl_get_explicit_stride(type->type));
      }

      case vtn_base_type_struct: {
         bool need_new_struct = false;
         const uint32_t num_fields = type->length;
         NIR_VLA(struct glsl_struct_field, fields, num_fields);
         for (unsigned i = 0; i < num_fields; i++) {
            fields[i] = *glsl_get_struct_field_data(type->type, i);
            const struct glsl_type *field_nir_type =
               vtn_type_get_nir_type(b, type->members[i], mode);
            if (fields[i].type != field_nir_type) {
               fields[i].type = field_nir_type;
               need_new_struct = true;
            }
         }
         if (need_new_struct) {
            if (glsl_type_is_interface(type->type)) {
               return glsl_interface_type(fields, num_fields,
                                          /* packing */ 0, false,
                                          glsl_get_type_name(type->type));
            } else {
               return glsl_struct_type(fields, num_fields,
                                       glsl_get_type_name(type->type),
                                       glsl_struct_type_is_packed(type->type));
            }
         } else {
            /* No changes, just pass it on */
            return type->type;
         }
      }

      case vtn_base_type_image:
         return type->glsl_image;

      case vtn_base_type_sampler:
         return glsl_bare_sampler_type();

      case vtn_base_type_sampled_image:
         return type->image->glsl_image;

      default:
         return type->type;
      }
   }

   /* Layout decorations are allowed but ignored in certain conditions,
    * to allow SPIR-V generators perform type deduplication.  Discard
    * unnecessary ones when passing to NIR.
    */
   if (!vtn_type_needs_explicit_layout(b, mode))
      return glsl_get_bare_type(type->type);

   return type->type;
}

static struct vtn_type *
mutable_matrix_member(struct vtn_builder *b, struct vtn_type *type, int member)
{
   type->members[member] = vtn_type_copy(b, type->members[member]);
   type = type->members[member];

   /* We may have an array of matrices.... Oh, joy! */
   while (glsl_type_is_array(type->type)) {
      type->array_element = vtn_type_copy(b, type->array_element);
      type = type->array_element;
   }

   vtn_assert(glsl_type_is_matrix(type->type));

   return type;
}

static void
vtn_handle_access_qualifier(struct vtn_builder *b, struct vtn_type *type,
                            int member, enum gl_access_qualifier access)
{
   type->members[member] = vtn_type_copy(b, type->members[member]);
   type = type->members[member];

   type->access |= access;
}

static void
array_stride_decoration_cb(struct vtn_builder *b,
                           struct vtn_value *val, int member,
                           const struct vtn_decoration *dec, void *void_ctx)
{
   struct vtn_type *type = val->type;

   if (dec->decoration == SpvDecorationArrayStride) {
      if (vtn_type_contains_block(b, type)) {
         vtn_warn("The ArrayStride decoration cannot be applied to an array "
                  "type which contains a structure type decorated Block "
                  "or BufferBlock");
         /* Ignore the decoration */
      } else {
         vtn_fail_if(dec->operands[0] == 0, "ArrayStride must be non-zero");
         type->stride = dec->operands[0];
      }
   }
}

static void
struct_member_decoration_cb(struct vtn_builder *b,
                            UNUSED struct vtn_value *val, int member,
                            const struct vtn_decoration *dec, void *void_ctx)
{
   struct member_decoration_ctx *ctx = void_ctx;

   if (member < 0)
      return;

   assert(member < ctx->num_fields);

   switch (dec->decoration) {
   case SpvDecorationRelaxedPrecision:
   case SpvDecorationUniform:
   case SpvDecorationUniformId:
      break; /* FIXME: Do nothing with this for now. */
   case SpvDecorationNonWritable:
      vtn_handle_access_qualifier(b, ctx->type, member, ACCESS_NON_WRITEABLE);
      break;
   case SpvDecorationNonReadable:
      vtn_handle_access_qualifier(b, ctx->type, member, ACCESS_NON_READABLE);
      break;
   case SpvDecorationVolatile:
      vtn_handle_access_qualifier(b, ctx->type, member, ACCESS_VOLATILE);
      break;
   case SpvDecorationCoherent:
      vtn_handle_access_qualifier(b, ctx->type, member, ACCESS_COHERENT);
      break;
   case SpvDecorationNoPerspective:
      ctx->fields[member].interpolation = INTERP_MODE_NOPERSPECTIVE;
      break;
   case SpvDecorationFlat:
      ctx->fields[member].interpolation = INTERP_MODE_FLAT;
      break;
   case SpvDecorationExplicitInterpAMD:
      ctx->fields[member].interpolation = INTERP_MODE_EXPLICIT;
      break;
   case SpvDecorationCentroid:
      ctx->fields[member].centroid = true;
      break;
   case SpvDecorationSample:
      ctx->fields[member].sample = true;
      break;
   case SpvDecorationStream:
      /* This is handled later by var_decoration_cb in vtn_variables.c */
      break;
   case SpvDecorationLocation:
      ctx->fields[member].location = dec->operands[0];
      break;
   case SpvDecorationComponent:
      break; /* FIXME: What should we do with these? */
   case SpvDecorationBuiltIn:
      ctx->type->members[member] = vtn_type_copy(b, ctx->type->members[member]);
      ctx->type->members[member]->is_builtin = true;
      ctx->type->members[member]->builtin = dec->operands[0];
      ctx->type->builtin_block = true;
      break;
   case SpvDecorationOffset:
      ctx->type->offsets[member] = dec->operands[0];
      ctx->fields[member].offset = dec->operands[0];
      break;
   case SpvDecorationMatrixStride:
      /* Handled as a second pass */
      break;
   case SpvDecorationColMajor:
      break; /* Nothing to do here.  Column-major is the default. */
   case SpvDecorationRowMajor:
      mutable_matrix_member(b, ctx->type, member)->row_major = true;
      break;

   case SpvDecorationPatch:
      break;

   case SpvDecorationSpecId:
   case SpvDecorationBlock:
   case SpvDecorationBufferBlock:
   case SpvDecorationArrayStride:
   case SpvDecorationGLSLShared:
   case SpvDecorationGLSLPacked:
   case SpvDecorationInvariant:
   case SpvDecorationRestrict:
   case SpvDecorationAliased:
   case SpvDecorationConstant:
   case SpvDecorationIndex:
   case SpvDecorationBinding:
   case SpvDecorationDescriptorSet:
   case SpvDecorationLinkageAttributes:
   case SpvDecorationNoContraction:
   case SpvDecorationInputAttachmentIndex:
      vtn_warn("Decoration not allowed on struct members: %s",
               spirv_decoration_to_string(dec->decoration));
      break;

   case SpvDecorationXfbBuffer:
   case SpvDecorationXfbStride:
      /* This is handled later by var_decoration_cb in vtn_variables.c */
      break;

   case SpvDecorationCPacked:
      if (b->shader->info.stage != MESA_SHADER_KERNEL)
         vtn_warn("Decoration only allowed for CL-style kernels: %s",
                  spirv_decoration_to_string(dec->decoration));
      else
         ctx->type->packed = true;
      break;

   case SpvDecorationSaturatedConversion:
   case SpvDecorationFuncParamAttr:
   case SpvDecorationFPRoundingMode:
   case SpvDecorationFPFastMathMode:
   case SpvDecorationAlignment:
      if (b->shader->info.stage != MESA_SHADER_KERNEL) {
         vtn_warn("Decoration only allowed for CL-style kernels: %s",
                  spirv_decoration_to_string(dec->decoration));
      }
      break;

   case SpvDecorationUserSemantic:
   case SpvDecorationUserTypeGOOGLE:
      /* User semantic decorations can safely be ignored by the driver. */
      break;

   default:
      vtn_fail_with_decoration("Unhandled decoration", dec->decoration);
   }
}

/** Chases the array type all the way down to the tail and rewrites the
 * glsl_types to be based off the tail's glsl_type.
 */
static void
vtn_array_type_rewrite_glsl_type(struct vtn_type *type)
{
   if (type->base_type != vtn_base_type_array)
      return;

   vtn_array_type_rewrite_glsl_type(type->array_element);

   type->type = glsl_array_type(type->array_element->type,
                                type->length, type->stride);
}

/* Matrix strides are handled as a separate pass because we need to know
 * whether the matrix is row-major or not first.
 */
static void
struct_member_matrix_stride_cb(struct vtn_builder *b,
                               UNUSED struct vtn_value *val, int member,
                               const struct vtn_decoration *dec,
                               void *void_ctx)
{
   if (dec->decoration != SpvDecorationMatrixStride)
      return;

   vtn_fail_if(member < 0,
               "The MatrixStride decoration is only allowed on members "
               "of OpTypeStruct");
   vtn_fail_if(dec->operands[0] == 0, "MatrixStride must be non-zero");

   struct member_decoration_ctx *ctx = void_ctx;

   struct vtn_type *mat_type = mutable_matrix_member(b, ctx->type, member);
   if (mat_type->row_major) {
      mat_type->array_element = vtn_type_copy(b, mat_type->array_element);
      mat_type->stride = mat_type->array_element->stride;
      mat_type->array_element->stride = dec->operands[0];

      mat_type->type = glsl_explicit_matrix_type(mat_type->type,
                                                 dec->operands[0], true);
      mat_type->array_element->type = glsl_get_column_type(mat_type->type);
   } else {
      vtn_assert(mat_type->array_element->stride > 0);
      mat_type->stride = dec->operands[0];

      mat_type->type = glsl_explicit_matrix_type(mat_type->type,
                                                 dec->operands[0], false);
   }

   /* Now that we've replaced the glsl_type with a properly strided matrix
    * type, rewrite the member type so that it's an array of the proper kind
    * of glsl_type.
    */
   vtn_array_type_rewrite_glsl_type(ctx->type->members[member]);
   ctx->fields[member].type = ctx->type->members[member]->type;
}

static void
struct_block_decoration_cb(struct vtn_builder *b,
                           struct vtn_value *val, int member,
                           const struct vtn_decoration *dec, void *ctx)
{
   if (member != -1)
      return;

   struct vtn_type *type = val->type;
   if (dec->decoration == SpvDecorationBlock)
      type->block = true;
   else if (dec->decoration == SpvDecorationBufferBlock)
      type->buffer_block = true;
}

static void
type_decoration_cb(struct vtn_builder *b,
                   struct vtn_value *val, int member,
                   const struct vtn_decoration *dec, UNUSED void *ctx)
{
   struct vtn_type *type = val->type;

   if (member != -1) {
      /* This should have been handled by OpTypeStruct */
      assert(val->type->base_type == vtn_base_type_struct);
      assert(member >= 0 && member < val->type->length);
      return;
   }

   switch (dec->decoration) {
   case SpvDecorationArrayStride:
      vtn_assert(type->base_type == vtn_base_type_array ||
                 type->base_type == vtn_base_type_pointer);
      break;
   case SpvDecorationBlock:
      vtn_assert(type->base_type == vtn_base_type_struct);
      vtn_assert(type->block);
      break;
   case SpvDecorationBufferBlock:
      vtn_assert(type->base_type == vtn_base_type_struct);
      vtn_assert(type->buffer_block);
      break;
   case SpvDecorationGLSLShared:
   case SpvDecorationGLSLPacked:
      /* Ignore these, since we get explicit offsets anyways */
      break;

   case SpvDecorationRowMajor:
   case SpvDecorationColMajor:
   case SpvDecorationMatrixStride:
   case SpvDecorationBuiltIn:
   case SpvDecorationNoPerspective:
   case SpvDecorationFlat:
   case SpvDecorationPatch:
   case SpvDecorationCentroid:
   case SpvDecorationSample:
   case SpvDecorationExplicitInterpAMD:
   case SpvDecorationVolatile:
   case SpvDecorationCoherent:
   case SpvDecorationNonWritable:
   case SpvDecorationNonReadable:
   case SpvDecorationUniform:
   case SpvDecorationUniformId:
   case SpvDecorationLocation:
   case SpvDecorationComponent:
   case SpvDecorationOffset:
   case SpvDecorationXfbBuffer:
   case SpvDecorationXfbStride:
   case SpvDecorationUserSemantic:
      vtn_warn("Decoration only allowed for struct members: %s",
               spirv_decoration_to_string(dec->decoration));
      break;

   case SpvDecorationStream:
      /* We don't need to do anything here, as stream is filled up when
       * aplying the decoration to a variable, just check that if it is not a
       * struct member, it should be a struct.
       */
      vtn_assert(type->base_type == vtn_base_type_struct);
      break;

   case SpvDecorationRelaxedPrecision:
   case SpvDecorationSpecId:
   case SpvDecorationInvariant:
   case SpvDecorationRestrict:
   case SpvDecorationAliased:
   case SpvDecorationConstant:
   case SpvDecorationIndex:
   case SpvDecorationBinding:
   case SpvDecorationDescriptorSet:
   case SpvDecorationLinkageAttributes:
   case SpvDecorationNoContraction:
   case SpvDecorationInputAttachmentIndex:
      vtn_warn("Decoration not allowed on types: %s",
               spirv_decoration_to_string(dec->decoration));
      break;

   case SpvDecorationCPacked:
      if (b->shader->info.stage != MESA_SHADER_KERNEL)
         vtn_warn("Decoration only allowed for CL-style kernels: %s",
                  spirv_decoration_to_string(dec->decoration));
      else
         type->packed = true;
      break;

   case SpvDecorationSaturatedConversion:
   case SpvDecorationFuncParamAttr:
   case SpvDecorationFPRoundingMode:
   case SpvDecorationFPFastMathMode:
   case SpvDecorationAlignment:
      vtn_warn("Decoration only allowed for CL-style kernels: %s",
               spirv_decoration_to_string(dec->decoration));
      break;

   case SpvDecorationUserTypeGOOGLE:
      /* User semantic decorations can safely be ignored by the driver. */
      break;

   default:
      vtn_fail_with_decoration("Unhandled decoration", dec->decoration);
   }
}

static unsigned
translate_image_format(struct vtn_builder *b, SpvImageFormat format)
{
   switch (format) {
   case SpvImageFormatUnknown:      return PIPE_FORMAT_NONE;
   case SpvImageFormatRgba32f:      return PIPE_FORMAT_R32G32B32A32_FLOAT;
   case SpvImageFormatRgba16f:      return PIPE_FORMAT_R16G16B16A16_FLOAT;
   case SpvImageFormatR32f:         return PIPE_FORMAT_R32_FLOAT;
   case SpvImageFormatRgba8:        return PIPE_FORMAT_R8G8B8A8_UNORM;
   case SpvImageFormatRgba8Snorm:   return PIPE_FORMAT_R8G8B8A8_SNORM;
   case SpvImageFormatRg32f:        return PIPE_FORMAT_R32G32_FLOAT;
   case SpvImageFormatRg16f:        return PIPE_FORMAT_R16G16_FLOAT;
   case SpvImageFormatR11fG11fB10f: return PIPE_FORMAT_R11G11B10_FLOAT;
   case SpvImageFormatR16f:         return PIPE_FORMAT_R16_FLOAT;
   case SpvImageFormatRgba16:       return PIPE_FORMAT_R16G16B16A16_UNORM;
   case SpvImageFormatRgb10A2:      return PIPE_FORMAT_R10G10B10A2_UNORM;
   case SpvImageFormatRg16:         return PIPE_FORMAT_R16G16_UNORM;
   case SpvImageFormatRg8:          return PIPE_FORMAT_R8G8_UNORM;
   case SpvImageFormatR16:          return PIPE_FORMAT_R16_UNORM;
   case SpvImageFormatR8:           return PIPE_FORMAT_R8_UNORM;
   case SpvImageFormatRgba16Snorm:  return PIPE_FORMAT_R16G16B16A16_SNORM;
   case SpvImageFormatRg16Snorm:    return PIPE_FORMAT_R16G16_SNORM;
   case SpvImageFormatRg8Snorm:     return PIPE_FORMAT_R8G8_SNORM;
   case SpvImageFormatR16Snorm:     return PIPE_FORMAT_R16_SNORM;
   case SpvImageFormatR8Snorm:      return PIPE_FORMAT_R8_SNORM;
   case SpvImageFormatRgba32i:      return PIPE_FORMAT_R32G32B32A32_SINT;
   case SpvImageFormatRgba16i:      return PIPE_FORMAT_R16G16B16A16_SINT;
   case SpvImageFormatRgba8i:       return PIPE_FORMAT_R8G8B8A8_SINT;
   case SpvImageFormatR32i:         return PIPE_FORMAT_R32_SINT;
   case SpvImageFormatRg32i:        return PIPE_FORMAT_R32G32_SINT;
   case SpvImageFormatRg16i:        return PIPE_FORMAT_R16G16_SINT;
   case SpvImageFormatRg8i:         return PIPE_FORMAT_R8G8_SINT;
   case SpvImageFormatR16i:         return PIPE_FORMAT_R16_SINT;
   case SpvImageFormatR8i:          return PIPE_FORMAT_R8_SINT;
   case SpvImageFormatRgba32ui:     return PIPE_FORMAT_R32G32B32A32_UINT;
   case SpvImageFormatRgba16ui:     return PIPE_FORMAT_R16G16B16A16_UINT;
   case SpvImageFormatRgba8ui:      return PIPE_FORMAT_R8G8B8A8_UINT;
   case SpvImageFormatR32ui:        return PIPE_FORMAT_R32_UINT;
   case SpvImageFormatRgb10a2ui:    return PIPE_FORMAT_R10G10B10A2_UINT;
   case SpvImageFormatRg32ui:       return PIPE_FORMAT_R32G32_UINT;
   case SpvImageFormatRg16ui:       return PIPE_FORMAT_R16G16_UINT;
   case SpvImageFormatRg8ui:        return PIPE_FORMAT_R8G8_UINT;
   case SpvImageFormatR16ui:        return PIPE_FORMAT_R16_UINT;
   case SpvImageFormatR8ui:         return PIPE_FORMAT_R8_UINT;
   default:
      vtn_fail("Invalid image format: %s (%u)",
               spirv_imageformat_to_string(format), format);
   }
}

static void
vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
                const uint32_t *w, unsigned count)
{
   struct vtn_value *val = NULL;

   /* In order to properly handle forward declarations, we have to defer
    * allocation for pointer types.
    */
   if (opcode != SpvOpTypePointer && opcode != SpvOpTypeForwardPointer) {
      val = vtn_push_value(b, w[1], vtn_value_type_type);
      vtn_fail_if(val->type != NULL,
                  "Only pointers can have forward declarations");
      val->type = rzalloc(b, struct vtn_type);
      val->type->id = w[1];
   }

   switch (opcode) {
   case SpvOpTypeVoid:
      val->type->base_type = vtn_base_type_void;
      val->type->type = glsl_void_type();
      break;
   case SpvOpTypeBool:
      val->type->base_type = vtn_base_type_scalar;
      val->type->type = glsl_bool_type();
      val->type->length = 1;
      break;
   case SpvOpTypeInt: {
      int bit_size = w[2];
      const bool signedness = w[3];
      val->type->base_type = vtn_base_type_scalar;
      switch (bit_size) {
      case 64:
         val->type->type = (signedness ? glsl_int64_t_type() : glsl_uint64_t_type());
         break;
      case 32:
         val->type->type = (signedness ? glsl_int_type() : glsl_uint_type());
         break;
      case 16:
         val->type->type = (signedness ? glsl_int16_t_type() : glsl_uint16_t_type());
         break;
      case 8:
         val->type->type = (signedness ? glsl_int8_t_type() : glsl_uint8_t_type());
         break;
      default:
         vtn_fail("Invalid int bit size: %u", bit_size);
      }
      val->type->length = 1;
      break;
   }

   case SpvOpTypeFloat: {
      int bit_size = w[2];
      val->type->base_type = vtn_base_type_scalar;
      switch (bit_size) {
      case 16:
         val->type->type = glsl_float16_t_type();
         break;
      case 32:
         val->type->type = glsl_float_type();
         break;
      case 64:
         val->type->type = glsl_double_type();
         break;
      default:
         vtn_fail("Invalid float bit size: %u", bit_size);
      }
      val->type->length = 1;
      break;
   }

   case SpvOpTypeVector: {
      struct vtn_type *base = vtn_get_type(b, w[2]);
      unsigned elems = w[3];

      vtn_fail_if(base->base_type != vtn_base_type_scalar,
                  "Base type for OpTypeVector must be a scalar");
      vtn_fail_if((elems < 2 || elems > 4) && (elems != 8) && (elems != 16),
                  "Invalid component count for OpTypeVector");

      val->type->base_type = vtn_base_type_vector;
      val->type->type = glsl_vector_type(glsl_get_base_type(base->type), elems);
      val->type->length = elems;
      val->type->stride = glsl_type_is_boolean(val->type->type)
         ? 4 : glsl_get_bit_size(base->type) / 8;
      val->type->array_element = base;
      break;
   }

   case SpvOpTypeMatrix: {
      struct vtn_type *base = vtn_get_type(b, w[2]);
      unsigned columns = w[3];

      vtn_fail_if(base->base_type != vtn_base_type_vector,
                  "Base type for OpTypeMatrix must be a vector");
      vtn_fail_if(columns < 2 || columns > 4,
                  "Invalid column count for OpTypeMatrix");

      val->type->base_type = vtn_base_type_matrix;
      val->type->type = glsl_matrix_type(glsl_get_base_type(base->type),
                                         glsl_get_vector_elements(base->type),
                                         columns);
      vtn_fail_if(glsl_type_is_error(val->type->type),
                  "Unsupported base type for OpTypeMatrix");
      assert(!glsl_type_is_error(val->type->type));
      val->type->length = columns;
      val->type->array_element = base;
      val->type->row_major = false;
      val->type->stride = 0;
      break;
   }

   case SpvOpTypeRuntimeArray:
   case SpvOpTypeArray: {
      struct vtn_type *array_element = vtn_get_type(b, w[2]);

      if (opcode == SpvOpTypeRuntimeArray) {
         /* A length of 0 is used to denote unsized arrays */
         val->type->length = 0;
      } else {
         val->type->length = vtn_constant_uint(b, w[3]);
      }

      val->type->base_type = vtn_base_type_array;
      val->type->array_element = array_element;
      if (b->shader->info.stage == MESA_SHADER_KERNEL)
         val->type->stride = glsl_get_cl_size(array_element->type);

      vtn_foreach_decoration(b, val, array_stride_decoration_cb, NULL);
      val->type->type = glsl_array_type(array_element->type, val->type->length,
                                        val->type->stride);
      break;
   }

   case SpvOpTypeStruct: {
      unsigned num_fields = count - 2;
      val->type->base_type = vtn_base_type_struct;
      val->type->length = num_fields;
      val->type->members = ralloc_array(b, struct vtn_type *, num_fields);
      val->type->offsets = ralloc_array(b, unsigned, num_fields);
      val->type->packed = false;

      NIR_VLA(struct glsl_struct_field, fields, count);
      for (unsigned i = 0; i < num_fields; i++) {
         val->type->members[i] = vtn_get_type(b, w[i + 2]);
         fields[i] = (struct glsl_struct_field) {
            .type = val->type->members[i]->type,
            .name = ralloc_asprintf(b, "field%d", i),
            .location = -1,
            .offset = -1,
         };
      }

      if (b->shader->info.stage == MESA_SHADER_KERNEL) {
         unsigned offset = 0;
         for (unsigned i = 0; i < num_fields; i++) {
            offset = align(offset, glsl_get_cl_alignment(fields[i].type));
            fields[i].offset = offset;
            offset += glsl_get_cl_size(fields[i].type);
         }
      }

      struct member_decoration_ctx ctx = {
         .num_fields = num_fields,
         .fields = fields,
         .type = val->type
      };

      vtn_foreach_decoration(b, val, struct_member_decoration_cb, &ctx);
      vtn_foreach_decoration(b, val, struct_member_matrix_stride_cb, &ctx);

      vtn_foreach_decoration(b, val, struct_block_decoration_cb, NULL);

      const char *name = val->name;

      if (val->type->block || val->type->buffer_block) {
         /* Packing will be ignored since types coming from SPIR-V are
          * explicitly laid out.
          */
         val->type->type = glsl_interface_type(fields, num_fields,
                                               /* packing */ 0, false,
                                               name ? name : "block");
      } else {
         val->type->type = glsl_struct_type(fields, num_fields,
                                            name ? name : "struct", false);
      }
      break;
   }

   case SpvOpTypeFunction: {
      val->type->base_type = vtn_base_type_function;
      val->type->type = NULL;

      val->type->return_type = vtn_get_type(b, w[2]);

      const unsigned num_params = count - 3;
      val->type->length = num_params;
      val->type->params = ralloc_array(b, struct vtn_type *, num_params);
      for (unsigned i = 0; i < count - 3; i++) {
         val->type->params[i] = vtn_get_type(b, w[i + 3]);
      }
      break;
   }

   case SpvOpTypePointer:
   case SpvOpTypeForwardPointer: {
      /* We can't blindly push the value because it might be a forward
       * declaration.
       */
      val = vtn_untyped_value(b, w[1]);

      SpvStorageClass storage_class = w[2];

      if (val->value_type == vtn_value_type_invalid) {
         val->value_type = vtn_value_type_type;
         val->type = rzalloc(b, struct vtn_type);
         val->type->id = w[1];
         val->type->base_type = vtn_base_type_pointer;
         val->type->storage_class = storage_class;

         /* These can actually be stored to nir_variables and used as SSA
          * values so they need a real glsl_type.
          */
         enum vtn_variable_mode mode = vtn_storage_class_to_mode(
            b, storage_class, NULL, NULL);
         val->type->type = nir_address_format_to_glsl_type(
            vtn_mode_to_address_format(b, mode));
      } else {
         vtn_fail_if(val->type->storage_class != storage_class,
                     "The storage classes of an OpTypePointer and any "
                     "OpTypeForwardPointers that provide forward "
                     "declarations of it must match.");
      }

      if (opcode == SpvOpTypePointer) {
         vtn_fail_if(val->type->deref != NULL,
                     "While OpTypeForwardPointer can be used to provide a "
                     "forward declaration of a pointer, OpTypePointer can "
                     "only be used once for a given id.");

         val->type->deref = vtn_get_type(b, w[3]);

         /* Only certain storage classes use ArrayStride.  The others (in
          * particular Workgroup) are expected to be laid out by the driver.
          */
         switch (storage_class) {
         case SpvStorageClassUniform:
         case SpvStorageClassPushConstant:
         case SpvStorageClassStorageBuffer:
         case SpvStorageClassPhysicalStorageBuffer:
            vtn_foreach_decoration(b, val, array_stride_decoration_cb, NULL);
            break;
         default:
            /* Nothing to do. */
            break;
         }

         if (b->physical_ptrs) {
            switch (storage_class) {
            case SpvStorageClassFunction:
            case SpvStorageClassWorkgroup:
            case SpvStorageClassCrossWorkgroup:
            case SpvStorageClassUniformConstant:
               val->type->stride = align(glsl_get_cl_size(val->type->deref->type),
                                         glsl_get_cl_alignment(val->type->deref->type));
               break;
            default:
               break;
            }
         }
      }
      break;
   }

   case SpvOpTypeImage: {
      val->type->base_type = vtn_base_type_image;

      /* Images are represented in NIR as a scalar SSA value that is the
       * result of a deref instruction.  An OpLoad on an OpTypeImage pointer
       * from UniformConstant memory just takes the NIR deref from the pointer
       * and turns it into an SSA value.
       */
      val->type->type = nir_address_format_to_glsl_type(
         vtn_mode_to_address_format(b, vtn_variable_mode_function));

      const struct vtn_type *sampled_type = vtn_get_type(b, w[2]);
      vtn_fail_if(sampled_type->base_type != vtn_base_type_scalar ||
                  glsl_get_bit_size(sampled_type->type) != 32,
                  "Sampled type of OpTypeImage must be a 32-bit scalar");

      enum glsl_sampler_dim dim;
      switch ((SpvDim)w[3]) {
      case SpvDim1D:       dim = GLSL_SAMPLER_DIM_1D;    break;
      case SpvDim2D:       dim = GLSL_SAMPLER_DIM_2D;    break;
      case SpvDim3D:       dim = GLSL_SAMPLER_DIM_3D;    break;
      case SpvDimCube:     dim = GLSL_SAMPLER_DIM_CUBE;  break;
      case SpvDimRect:     dim = GLSL_SAMPLER_DIM_RECT;  break;
      case SpvDimBuffer:   dim = GLSL_SAMPLER_DIM_BUF;   break;
      case SpvDimSubpassData: dim = GLSL_SAMPLER_DIM_SUBPASS; break;
      default:
         vtn_fail("Invalid SPIR-V image dimensionality: %s (%u)",
                  spirv_dim_to_string((SpvDim)w[3]), w[3]);
      }

      /* w[4]: as per Vulkan spec "Validation Rules within a Module",
       *       The “Depth” operand of OpTypeImage is ignored.
       */
      bool is_array = w[5];
      bool multisampled = w[6];
      unsigned sampled = w[7];
      SpvImageFormat format = w[8];

      if (count > 9)
         val->type->access_qualifier = w[9];
      else
         val->type->access_qualifier = SpvAccessQualifierReadWrite;

      if (multisampled) {
         if (dim == GLSL_SAMPLER_DIM_2D)
            dim = GLSL_SAMPLER_DIM_MS;
         else if (dim == GLSL_SAMPLER_DIM_SUBPASS)
            dim = GLSL_SAMPLER_DIM_SUBPASS_MS;
         else
            vtn_fail("Unsupported multisampled image type");
      }

      val->type->image_format = translate_image_format(b, format);

      enum glsl_base_type sampled_base_type =
         glsl_get_base_type(sampled_type->type);
      if (sampled == 1) {
         val->type->glsl_image = glsl_sampler_type(dim, false, is_array,
                                                   sampled_base_type);
      } else if (sampled == 2) {
         val->type->glsl_image = glsl_image_type(dim, is_array,
                                                 sampled_base_type);
      } else {
         vtn_fail("We need to know if the image will be sampled");
      }
      break;
   }

   case SpvOpTypeSampledImage: {
      val->type->base_type = vtn_base_type_sampled_image;
      val->type->image = vtn_get_type(b, w[2]);

      /* Sampled images are represented NIR as a vec2 SSA value where each
       * component is the result of a deref instruction.  The first component
       * is the image and the second is the sampler.  An OpLoad on an
       * OpTypeSampledImage pointer from UniformConstant memory just takes
       * the NIR deref from the pointer and duplicates it to both vector
       * components.
       */
      nir_address_format addr_format =
         vtn_mode_to_address_format(b, vtn_variable_mode_function);
      assert(nir_address_format_num_components(addr_format) == 1);
      unsigned bit_size = nir_address_format_bit_size(addr_format);
      assert(bit_size == 32 || bit_size == 64);

      enum glsl_base_type base_type =
         bit_size == 32 ? GLSL_TYPE_UINT : GLSL_TYPE_UINT64;
      val->type->type = glsl_vector_type(base_type, 2);
      break;
   }

   case SpvOpTypeSampler:
      val->type->base_type = vtn_base_type_sampler;

      /* Samplers are represented in NIR as a scalar SSA value that is the
       * result of a deref instruction.  An OpLoad on an OpTypeSampler pointer
       * from UniformConstant memory just takes the NIR deref from the pointer
       * and turns it into an SSA value.
       */
      val->type->type = nir_address_format_to_glsl_type(
         vtn_mode_to_address_format(b, vtn_variable_mode_function));
      break;

   case SpvOpTypeOpaque:
   case SpvOpTypeEvent:
   case SpvOpTypeDeviceEvent:
   case SpvOpTypeReserveId:
   case SpvOpTypeQueue:
   case SpvOpTypePipe:
   default:
      vtn_fail_with_opcode("Unhandled opcode", opcode);
   }

   vtn_foreach_decoration(b, val, type_decoration_cb, NULL);

   if (val->type->base_type == vtn_base_type_struct &&
       (val->type->block || val->type->buffer_block)) {
      for (unsigned i = 0; i < val->type->length; i++) {
         vtn_fail_if(vtn_type_contains_block(b, val->type->members[i]),
                     "Block and BufferBlock decorations cannot decorate a "
                     "structure type that is nested at any level inside "
                     "another structure type decorated with Block or "
                     "BufferBlock.");
      }
   }
}

static nir_constant *
vtn_null_constant(struct vtn_builder *b, struct vtn_type *type)
{
   nir_constant *c = rzalloc(b, nir_constant);

   switch (type->base_type) {
   case vtn_base_type_scalar:
   case vtn_base_type_vector:
      /* Nothing to do here.  It's already initialized to zero */
      break;

   case vtn_base_type_pointer: {
      enum vtn_variable_mode mode = vtn_storage_class_to_mode(
         b, type->storage_class, type->deref, NULL);
      nir_address_format addr_format = vtn_mode_to_address_format(b, mode);

      const nir_const_value *null_value = nir_address_format_null_value(addr_format);
      memcpy(c->values, null_value,
             sizeof(nir_const_value) * nir_address_format_num_components(addr_format));
      break;
   }

   case vtn_base_type_void:
   case vtn_base_type_image:
   case vtn_base_type_sampler:
   case vtn_base_type_sampled_image:
   case vtn_base_type_function:
      /* For those we have to return something but it doesn't matter what. */
      break;

   case vtn_base_type_matrix:
   case vtn_base_type_array:
      vtn_assert(type->length > 0);
      c->num_elements = type->length;
      c->elements = ralloc_array(b, nir_constant *, c->num_elements);

      c->elements[0] = vtn_null_constant(b, type->array_element);
      for (unsigned i = 1; i < c->num_elements; i++)
         c->elements[i] = c->elements[0];
      break;

   case vtn_base_type_struct:
      c->num_elements = type->length;
      c->elements = ralloc_array(b, nir_constant *, c->num_elements);
      for (unsigned i = 0; i < c->num_elements; i++)
         c->elements[i] = vtn_null_constant(b, type->members[i]);
      break;

   default:
      vtn_fail("Invalid type for null constant");
   }

   return c;
}

static void
spec_constant_decoration_cb(struct vtn_builder *b, UNUSED struct vtn_value *val,
                            ASSERTED int member,
                            const struct vtn_decoration *dec, void *data)
{
   vtn_assert(member == -1);
   if (dec->decoration != SpvDecorationSpecId)
      return;

   nir_const_value *value = data;
   for (unsigned i = 0; i < b->num_specializations; i++) {
      if (b->specializations[i].id == dec->operands[0]) {
         *value = b->specializations[i].value;
         return;
      }
   }
}

static void
handle_workgroup_size_decoration_cb(struct vtn_builder *b,
                                    struct vtn_value *val,
                                    ASSERTED int member,
                                    const struct vtn_decoration *dec,
                                    UNUSED void *data)
{
   vtn_assert(member == -1);
   if (dec->decoration != SpvDecorationBuiltIn ||
       dec->operands[0] != SpvBuiltInWorkgroupSize)
      return;

   vtn_assert(val->type->type == glsl_vector_type(GLSL_TYPE_UINT, 3));
   b->workgroup_size_builtin = val;
}

static void
vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
                    const uint32_t *w, unsigned count)
{
   struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_constant);
   val->constant = rzalloc(b, nir_constant);
   switch (opcode) {
   case SpvOpConstantTrue:
   case SpvOpConstantFalse:
   case SpvOpSpecConstantTrue:
   case SpvOpSpecConstantFalse: {
      vtn_fail_if(val->type->type != glsl_bool_type(),
                  "Result type of %s must be OpTypeBool",
                  spirv_op_to_string(opcode));

      bool bval = (opcode == SpvOpConstantTrue ||
                   opcode == SpvOpSpecConstantTrue);

      nir_const_value u32val = nir_const_value_for_uint(bval, 32);

      if (opcode == SpvOpSpecConstantTrue ||
          opcode == SpvOpSpecConstantFalse)
         vtn_foreach_decoration(b, val, spec_constant_decoration_cb, &u32val);

      val->constant->values[0].b = u32val.u32 != 0;
      break;
   }

   case SpvOpConstant:
   case SpvOpSpecConstant: {
      vtn_fail_if(val->type->base_type != vtn_base_type_scalar,
                  "Result type of %s must be a scalar",
                  spirv_op_to_string(opcode));
      int bit_size = glsl_get_bit_size(val->type->type);
      switch (bit_size) {
      case 64:
         val->constant->values[0].u64 = vtn_u64_literal(&w[3]);
         break;
      case 32:
         val->constant->values[0].u32 = w[3];
         break;
      case 16:
         val->constant->values[0].u16 = w[3];
         break;
      case 8:
         val->constant->values[0].u8 = w[3];
         break;
      default:
         vtn_fail("Unsupported SpvOpConstant bit size: %u", bit_size);
      }

      if (opcode == SpvOpSpecConstant)
         vtn_foreach_decoration(b, val, spec_constant_decoration_cb,
                                &val->constant->values[0]);
      break;
   }

   case SpvOpSpecConstantComposite:
   case SpvOpConstantComposite: {
      unsigned elem_count = count - 3;
      vtn_fail_if(elem_count != val->type->length,
                  "%s has %u constituents, expected %u",
                  spirv_op_to_string(opcode), elem_count, val->type->length);

      nir_constant **elems = ralloc_array(b, nir_constant *, elem_count);
      for (unsigned i = 0; i < elem_count; i++) {
         struct vtn_value *val = vtn_untyped_value(b, w[i + 3]);

         if (val->value_type == vtn_value_type_constant) {
            elems[i] = val->constant;
         } else {
            vtn_fail_if(val->value_type != vtn_value_type_undef,
                        "only constants or undefs allowed for "
                        "SpvOpConstantComposite");
            /* to make it easier, just insert a NULL constant for now */
            elems[i] = vtn_null_constant(b, val->type);
         }
      }

      switch (val->type->base_type) {
      case vtn_base_type_vector: {
         assert(glsl_type_is_vector(val->type->type));
         for (unsigned i = 0; i < elem_count; i++)
            val->constant->values[i] = elems[i]->values[0];
         break;
      }

      case vtn_base_type_matrix:
      case vtn_base_type_struct:
      case vtn_base_type_array:
         ralloc_steal(val->constant, elems);
         val->constant->num_elements = elem_count;
         val->constant->elements = elems;
         break;

      default:
         vtn_fail("Result type of %s must be a composite type",
                  spirv_op_to_string(opcode));
      }
      break;
   }

   case SpvOpSpecConstantOp: {
      nir_const_value u32op = nir_const_value_for_uint(w[3], 32);
      vtn_foreach_decoration(b, val, spec_constant_decoration_cb, &u32op);
      SpvOp opcode = u32op.u32;
      switch (opcode) {
      case SpvOpVectorShuffle: {
         struct vtn_value *v0 = &b->values[w[4]];
         struct vtn_value *v1 = &b->values[w[5]];

         vtn_assert(v0->value_type == vtn_value_type_constant ||
                    v0->value_type == vtn_value_type_undef);
         vtn_assert(v1->value_type == vtn_value_type_constant ||
                    v1->value_type == vtn_value_type_undef);

         unsigned len0 = glsl_get_vector_elements(v0->type->type);
         unsigned len1 = glsl_get_vector_elements(v1->type->type);

         vtn_assert(len0 + len1 < 16);

         unsigned bit_size = glsl_get_bit_size(val->type->type);
         unsigned bit_size0 = glsl_get_bit_size(v0->type->type);
         unsigned bit_size1 = glsl_get_bit_size(v1->type->type);

         vtn_assert(bit_size == bit_size0 && bit_size == bit_size1);
         (void)bit_size0; (void)bit_size1;

         nir_const_value undef = { .u64 = 0xdeadbeefdeadbeef };
         nir_const_value combined[NIR_MAX_VEC_COMPONENTS * 2];

         if (v0->value_type == vtn_value_type_constant) {
            for (unsigned i = 0; i < len0; i++)
               combined[i] = v0->constant->values[i];
         }
         if (v1->value_type == vtn_value_type_constant) {
            for (unsigned i = 0; i < len1; i++)
               combined[len0 + i] = v1->constant->values[i];
         }

         for (unsigned i = 0, j = 0; i < count - 6; i++, j++) {
            uint32_t comp = w[i + 6];
            if (comp == (uint32_t)-1) {
               /* If component is not used, set the value to a known constant
                * to detect if it is wrongly used.
                */
               val->constant->values[j] = undef;
            } else {
               vtn_fail_if(comp >= len0 + len1,
                           "All Component literals must either be FFFFFFFF "
                           "or in [0, N - 1] (inclusive).");
               val->constant->values[j] = combined[comp];
            }
         }
         break;
      }

      case SpvOpCompositeExtract:
      case SpvOpCompositeInsert: {
         struct vtn_value *comp;
         unsigned deref_start;
         struct nir_constant **c;
         if (opcode == SpvOpCompositeExtract) {
            comp = vtn_value(b, w[4], vtn_value_type_constant);
            deref_start = 5;
            c = &comp->constant;
         } else {
            comp = vtn_value(b, w[5], vtn_value_type_constant);
            deref_start = 6;
            val->constant = nir_constant_clone(comp->constant,
                                               (nir_variable *)b);
            c = &val->constant;
         }

         int elem = -1;
         const struct vtn_type *type = comp->type;
         for (unsigned i = deref_start; i < count; i++) {
            vtn_fail_if(w[i] > type->length,
                        "%uth index of %s is %u but the type has only "
                        "%u elements", i - deref_start,
                        spirv_op_to_string(opcode), w[i], type->length);

            switch (type->base_type) {
            case vtn_base_type_vector:
               elem = w[i];
               type = type->array_element;
               break;

            case vtn_base_type_matrix:
            case vtn_base_type_array:
               c = &(*c)->elements[w[i]];
               type = type->array_element;
               break;

            case vtn_base_type_struct:
               c = &(*c)->elements[w[i]];
               type = type->members[w[i]];
               break;

            default:
               vtn_fail("%s must only index into composite types",
                        spirv_op_to_string(opcode));
            }
         }

         if (opcode == SpvOpCompositeExtract) {
            if (elem == -1) {
               val->constant = *c;
            } else {
               unsigned num_components = type->length;
               for (unsigned i = 0; i < num_components; i++)
                  val->constant->values[i] = (*c)->values[elem + i];
            }
         } else {
            struct vtn_value *insert =
               vtn_value(b, w[4], vtn_value_type_constant);
            vtn_assert(insert->type == type);
            if (elem == -1) {
               *c = insert->constant;
            } else {
               unsigned num_components = type->length;
               for (unsigned i = 0; i < num_components; i++)
                  (*c)->values[elem + i] = insert->constant->values[i];
            }
         }
         break;
      }

      default: {
         bool swap;
         nir_alu_type dst_alu_type = nir_get_nir_type_for_glsl_type(val->type->type);
         nir_alu_type src_alu_type = dst_alu_type;
         unsigned num_components = glsl_get_vector_elements(val->type->type);
         unsigned bit_size;

         vtn_assert(count <= 7);

         switch (opcode) {
         case SpvOpSConvert:
         case SpvOpFConvert:
         case SpvOpUConvert:
            /* We have a source in a conversion */
            src_alu_type =
               nir_get_nir_type_for_glsl_type(vtn_get_value_type(b, w[4])->type);
            /* We use the bitsize of the conversion source to evaluate the opcode later */
            bit_size = glsl_get_bit_size(vtn_get_value_type(b, w[4])->type);
            break;
         default:
            bit_size = glsl_get_bit_size(val->type->type);
         };

         nir_op op = vtn_nir_alu_op_for_spirv_opcode(b, opcode, &swap,
                                                     nir_alu_type_get_type_size(src_alu_type),
                                                     nir_alu_type_get_type_size(dst_alu_type));
         nir_const_value src[3][NIR_MAX_VEC_COMPONENTS];

         for (unsigned i = 0; i < count - 4; i++) {
            struct vtn_value *src_val =
               vtn_value(b, w[4 + i], vtn_value_type_constant);

            /* If this is an unsized source, pull the bit size from the
             * source; otherwise, we'll use the bit size from the destination.
             */
            if (!nir_alu_type_get_type_size(nir_op_infos[op].input_types[i]))
               bit_size = glsl_get_bit_size(src_val->type->type);

            unsigned src_comps = nir_op_infos[op].input_sizes[i] ?
                                 nir_op_infos[op].input_sizes[i] :
                                 num_components;

            unsigned j = swap ? 1 - i : i;
            for (unsigned c = 0; c < src_comps; c++)
               src[j][c] = src_val->constant->values[c];
         }

         /* fix up fixed size sources */
         switch (op) {
         case nir_op_ishl:
         case nir_op_ishr:
         case nir_op_ushr: {
            if (bit_size == 32)
               break;
            for (unsigned i = 0; i < num_components; ++i) {
               switch (bit_size) {
               case 64: src[1][i].u32 = src[1][i].u64; break;
               case 16: src[1][i].u32 = src[1][i].u16; break;
               case  8: src[1][i].u32 = src[1][i].u8;  break;
               }
            }
            break;
         }
         default:
            break;
         }

         nir_const_value *srcs[3] = {
            src[0], src[1], src[2],
         };
         nir_eval_const_opcode(op, val->constant->values,
                               num_components, bit_size, srcs,
                               b->shader->info.float_controls_execution_mode);
         break;
      } /* default */
      }
      break;
   }

   case SpvOpConstantNull:
      val->constant = vtn_null_constant(b, val->type);
      break;

   case SpvOpConstantSampler:
      vtn_fail("OpConstantSampler requires Kernel Capability");
      break;

   default:
      vtn_fail_with_opcode("Unhandled opcode", opcode);
   }

   /* Now that we have the value, update the workgroup size if needed */
   vtn_foreach_decoration(b, val, handle_workgroup_size_decoration_cb, NULL);
}

SpvMemorySemanticsMask
vtn_storage_class_to_memory_semantics(SpvStorageClass sc)
{
   switch (sc) {
   case SpvStorageClassStorageBuffer:
   case SpvStorageClassPhysicalStorageBuffer:
      return SpvMemorySemanticsUniformMemoryMask;
   case SpvStorageClassWorkgroup:
      return SpvMemorySemanticsWorkgroupMemoryMask;
   default:
      return SpvMemorySemanticsMaskNone;
   }
}

static void
vtn_split_barrier_semantics(struct vtn_builder *b,
                            SpvMemorySemanticsMask semantics,
                            SpvMemorySemanticsMask *before,
                            SpvMemorySemanticsMask *after)
{
   /* For memory semantics embedded in operations, we split them into up to
    * two barriers, to be added before and after the operation.  This is less
    * strict than if we propagated until the final backend stage, but still
    * result in correct execution.
    *
    * A further improvement could be pipe this information (and use!) into the
    * next compiler layers, at the expense of making the handling of barriers
    * more complicated.
    */

   *before = SpvMemorySemanticsMaskNone;
   *after = SpvMemorySemanticsMaskNone;

   SpvMemorySemanticsMask order_semantics =
      semantics & (SpvMemorySemanticsAcquireMask |
                   SpvMemorySemanticsReleaseMask |
                   SpvMemorySemanticsAcquireReleaseMask |
                   SpvMemorySemanticsSequentiallyConsistentMask);

   if (util_bitcount(order_semantics) > 1) {
      /* Old GLSLang versions incorrectly set all the ordering bits.  This was
       * fixed in c51287d744fb6e7e9ccc09f6f8451e6c64b1dad6 of glslang repo,
       * and it is in GLSLang since revision "SPIRV99.1321" (from Jul-2016).
       */
      vtn_warn("Multiple memory ordering semantics specified, "
               "assuming AcquireRelease.");
      order_semantics = SpvMemorySemanticsAcquireReleaseMask;
   }

   const SpvMemorySemanticsMask av_vis_semantics =
      semantics & (SpvMemorySemanticsMakeAvailableMask |
                   SpvMemorySemanticsMakeVisibleMask);

   const SpvMemorySemanticsMask storage_semantics =
      semantics & (SpvMemorySemanticsUniformMemoryMask |
                   SpvMemorySemanticsSubgroupMemoryMask |
                   SpvMemorySemanticsWorkgroupMemoryMask |
                   SpvMemorySemanticsCrossWorkgroupMemoryMask |
                   SpvMemorySemanticsAtomicCounterMemoryMask |
                   SpvMemorySemanticsImageMemoryMask |
                   SpvMemorySemanticsOutputMemoryMask);

   const SpvMemorySemanticsMask other_semantics =
      semantics & ~(order_semantics | av_vis_semantics | storage_semantics);

   if (other_semantics)
      vtn_warn("Ignoring unhandled memory semantics: %u\n", other_semantics);

   /* SequentiallyConsistent is treated as AcquireRelease. */

   /* The RELEASE barrier happens BEFORE the operation, and it is usually
    * associated with a Store.  All the write operations with a matching
    * semantics will not be reordered after the Store.
    */
   if (order_semantics & (SpvMemorySemanticsReleaseMask |
                          SpvMemorySemanticsAcquireReleaseMask |
                          SpvMemorySemanticsSequentiallyConsistentMask)) {
      *before |= SpvMemorySemanticsReleaseMask | storage_semantics;
   }

   /* The ACQUIRE barrier happens AFTER the operation, and it is usually
    * associated with a Load.  All the operations with a matching semantics
    * will not be reordered before the Load.
    */
   if (order_semantics & (SpvMemorySemanticsAcquireMask |
                          SpvMemorySemanticsAcquireReleaseMask |
                          SpvMemorySemanticsSequentiallyConsistentMask)) {
      *after |= SpvMemorySemanticsAcquireMask | storage_semantics;
   }

   if (av_vis_semantics & SpvMemorySemanticsMakeVisibleMask)
      *before |= SpvMemorySemanticsMakeVisibleMask | storage_semantics;

   if (av_vis_semantics & SpvMemorySemanticsMakeAvailableMask)
      *after |= SpvMemorySemanticsMakeAvailableMask | storage_semantics;
}

static nir_memory_semantics
vtn_mem_semantics_to_nir_mem_semantics(struct vtn_builder *b,
                                       SpvMemorySemanticsMask semantics)
{
   nir_memory_semantics nir_semantics = 0;

   SpvMemorySemanticsMask order_semantics =
      semantics & (SpvMemorySemanticsAcquireMask |
                   SpvMemorySemanticsReleaseMask |
                   SpvMemorySemanticsAcquireReleaseMask |
                   SpvMemorySemanticsSequentiallyConsistentMask);

   if (util_bitcount(order_semantics) > 1) {
      /* Old GLSLang versions incorrectly set all the ordering bits.  This was
       * fixed in c51287d744fb6e7e9ccc09f6f8451e6c64b1dad6 of glslang repo,
       * and it is in GLSLang since revision "SPIRV99.1321" (from Jul-2016).
       */
      vtn_warn("Multiple memory ordering semantics bits specified, "
               "assuming AcquireRelease.");
      order_semantics = SpvMemorySemanticsAcquireReleaseMask;
   }

   switch (order_semantics) {
   case 0:
      /* Not an ordering barrier. */
      break;

   case SpvMemorySemanticsAcquireMask:
      nir_semantics = NIR_MEMORY_ACQUIRE;
      break;

   case SpvMemorySemanticsReleaseMask:
      nir_semantics = NIR_MEMORY_RELEASE;
      break;

   case SpvMemorySemanticsSequentiallyConsistentMask:
      /* Fall through.  Treated as AcquireRelease in Vulkan. */
   case SpvMemorySemanticsAcquireReleaseMask:
      nir_semantics = NIR_MEMORY_ACQUIRE | NIR_MEMORY_RELEASE;
      break;

   default:
      unreachable("Invalid memory order semantics");
   }

   if (semantics & SpvMemorySemanticsMakeAvailableMask) {
      vtn_fail_if(!b->options->caps.vk_memory_model,
                  "To use MakeAvailable memory semantics the VulkanMemoryModel "
                  "capability must be declared.");
      nir_semantics |= NIR_MEMORY_MAKE_AVAILABLE;
   }

   if (semantics & SpvMemorySemanticsMakeVisibleMask) {
      vtn_fail_if(!b->options->caps.vk_memory_model,
                  "To use MakeVisible memory semantics the VulkanMemoryModel "
                  "capability must be declared.");
      nir_semantics |= NIR_MEMORY_MAKE_VISIBLE;
   }

   return nir_semantics;
}

static nir_variable_mode
vtn_mem_sematics_to_nir_var_modes(struct vtn_builder *b,
                                  SpvMemorySemanticsMask semantics)
{
   /* Vulkan Environment for SPIR-V says "SubgroupMemory, CrossWorkgroupMemory,
    * and AtomicCounterMemory are ignored".
    */
   semantics &= ~(SpvMemorySemanticsSubgroupMemoryMask |
                  SpvMemorySemanticsCrossWorkgroupMemoryMask |
                  SpvMemorySemanticsAtomicCounterMemoryMask);

   /* TODO: Consider adding nir_var_mem_image mode to NIR so it can be used
    * for SpvMemorySemanticsImageMemoryMask.
    */

   nir_variable_mode modes = 0;
   if (semantics & (SpvMemorySemanticsUniformMemoryMask |
                    SpvMemorySemanticsImageMemoryMask)) {
      modes |= nir_var_uniform |
               nir_var_mem_ubo |
               nir_var_mem_ssbo |
               nir_var_mem_global;
   }
   if (semantics & SpvMemorySemanticsWorkgroupMemoryMask)
      modes |= nir_var_mem_shared;
   if (semantics & SpvMemorySemanticsOutputMemoryMask) {
      modes |= nir_var_shader_out;
   }

   return modes;
}

static nir_scope
vtn_scope_to_nir_scope(struct vtn_builder *b, SpvScope scope)
{
   nir_scope nir_scope;
   switch (scope) {
   case SpvScopeDevice:
      vtn_fail_if(b->options->caps.vk_memory_model &&
                  !b->options->caps.vk_memory_model_device_scope,
                  "If the Vulkan memory model is declared and any instruction "
                  "uses Device scope, the VulkanMemoryModelDeviceScope "
                  "capability must be declared.");
      nir_scope = NIR_SCOPE_DEVICE;
      break;

   case SpvScopeQueueFamily:
      vtn_fail_if(!b->options->caps.vk_memory_model,
                  "To use Queue Family scope, the VulkanMemoryModel capability "
                  "must be declared.");
      nir_scope = NIR_SCOPE_QUEUE_FAMILY;
      break;

   case SpvScopeWorkgroup:
      nir_scope = NIR_SCOPE_WORKGROUP;
      break;

   case SpvScopeSubgroup:
      nir_scope = NIR_SCOPE_SUBGROUP;
      break;

   case SpvScopeInvocation:
      nir_scope = NIR_SCOPE_INVOCATION;
      break;

   default:
      vtn_fail("Invalid memory scope");
   }

   return nir_scope;
}

static void
vtn_emit_scoped_control_barrier(struct vtn_builder *b, SpvScope exec_scope,
                                SpvScope mem_scope,
                                SpvMemorySemanticsMask semantics)
{
   nir_memory_semantics nir_semantics =
      vtn_mem_semantics_to_nir_mem_semantics(b, semantics);
   nir_variable_mode modes = vtn_mem_sematics_to_nir_var_modes(b, semantics);
   nir_scope nir_exec_scope = vtn_scope_to_nir_scope(b, exec_scope);

   /* Memory semantics is optional for OpControlBarrier. */
   nir_scope nir_mem_scope;
   if (nir_semantics == 0 || modes == 0)
      nir_mem_scope = NIR_SCOPE_NONE;
   else
      nir_mem_scope = vtn_scope_to_nir_scope(b, mem_scope);

   nir_scoped_barrier(&b->nb, nir_exec_scope, nir_mem_scope, nir_semantics, modes);
}

static void
vtn_emit_scoped_memory_barrier(struct vtn_builder *b, SpvScope scope,
                               SpvMemorySemanticsMask semantics)
{
   nir_variable_mode modes = vtn_mem_sematics_to_nir_var_modes(b, semantics);
   nir_memory_semantics nir_semantics =
      vtn_mem_semantics_to_nir_mem_semantics(b, semantics);

   /* No barrier to add. */
   if (nir_semantics == 0 || modes == 0)
      return;

   nir_scope nir_mem_scope = vtn_scope_to_nir_scope(b, scope);
   nir_scoped_barrier(&b->nb, NIR_SCOPE_NONE, nir_mem_scope, nir_semantics, modes);
}

struct vtn_ssa_value *
vtn_create_ssa_value(struct vtn_builder *b, const struct glsl_type *type)
{
   /* Always use bare types for SSA values for a couple of reasons:
    *
    *  1. Code which emits deref chains should never listen to the explicit
    *     layout information on the SSA value if any exists.  If we've
    *     accidentally been relying on this, we want to find those bugs.
    *
    *  2. We want to be able to quickly check that an SSA value being assigned
    *     to a SPIR-V value has the right type.  Using bare types everywhere
    *     ensures that we can pointer-compare.
    */
   struct vtn_ssa_value *val = rzalloc(b, struct vtn_ssa_value);
   val->type = glsl_get_bare_type(type);


   if (!glsl_type_is_vector_or_scalar(type)) {
      unsigned elems = glsl_get_length(val->type);
      val->elems = ralloc_array(b, struct vtn_ssa_value *, elems);
      if (glsl_type_is_array_or_matrix(type)) {
         const struct glsl_type *elem_type = glsl_get_array_element(type);
         for (unsigned i = 0; i < elems; i++)
            val->elems[i] = vtn_create_ssa_value(b, elem_type);
      } else {
         vtn_assert(glsl_type_is_struct_or_ifc(type));
         for (unsigned i = 0; i < elems; i++) {
            const struct glsl_type *elem_type = glsl_get_struct_field(type, i);
            val->elems[i] = vtn_create_ssa_value(b, elem_type);
         }
      }
   }

   return val;
}

static nir_tex_src
vtn_tex_src(struct vtn_builder *b, unsigned index, nir_tex_src_type type)
{
   nir_tex_src src;
   src.src = nir_src_for_ssa(vtn_get_nir_ssa(b, index));
   src.src_type = type;
   return src;
}

static uint32_t
image_operand_arg(struct vtn_builder *b, const uint32_t *w, uint32_t count,
                  uint32_t mask_idx, SpvImageOperandsMask op)
{
   static const SpvImageOperandsMask ops_with_arg =
      SpvImageOperandsBiasMask |
      SpvImageOperandsLodMask |
      SpvImageOperandsGradMask |
      SpvImageOperandsConstOffsetMask |
      SpvImageOperandsOffsetMask |
      SpvImageOperandsConstOffsetsMask |
      SpvImageOperandsSampleMask |
      SpvImageOperandsMinLodMask |
      SpvImageOperandsMakeTexelAvailableMask |
      SpvImageOperandsMakeTexelVisibleMask;

   assert(util_bitcount(op) == 1);
   assert(w[mask_idx] & op);
   assert(op & ops_with_arg);

   uint32_t idx = util_bitcount(w[mask_idx] & (op - 1) & ops_with_arg) + 1;

   /* Adjust indices for operands with two arguments. */
   static const SpvImageOperandsMask ops_with_two_args =
      SpvImageOperandsGradMask;
   idx += util_bitcount(w[mask_idx] & (op - 1) & ops_with_two_args);

   idx += mask_idx;

   vtn_fail_if(idx + (op & ops_with_two_args ? 1 : 0) >= count,
               "Image op claims to have %s but does not enough "
               "following operands", spirv_imageoperands_to_string(op));

   return idx;
}

static void
non_uniform_decoration_cb(struct vtn_builder *b,
                          struct vtn_value *val, int member,
                          const struct vtn_decoration *dec, void *void_ctx)
{
   enum gl_access_qualifier *access = void_ctx;
   switch (dec->decoration) {
   case SpvDecorationNonUniformEXT:
      *access |= ACCESS_NON_UNIFORM;
      break;

   default:
      break;
   }
}

static void
vtn_handle_texture(struct vtn_builder *b, SpvOp opcode,
                   const uint32_t *w, unsigned count)
{
   struct vtn_type *ret_type = vtn_get_type(b, w[1]);

   if (opcode == SpvOpSampledImage) {
      struct vtn_sampled_image si = {
         .image = vtn_get_image(b, w[3]),
         .sampler = vtn_get_sampler(b, w[4]),
      };
      vtn_push_sampled_image(b, w[2], si);
      return;
   } else if (opcode == SpvOpImage) {
      struct vtn_sampled_image si = vtn_get_sampled_image(b, w[3]);
      vtn_push_image(b, w[2], si.image);
      return;
   }

   nir_deref_instr *image = NULL, *sampler = NULL;
   struct vtn_value *sampled_val = vtn_untyped_value(b, w[3]);
   if (sampled_val->type->base_type == vtn_base_type_sampled_image) {
      struct vtn_sampled_image si = vtn_get_sampled_image(b, w[3]);
      image = si.image;
      sampler = si.sampler;
   } else {
      image = vtn_get_image(b, w[3]);
   }

   const enum glsl_sampler_dim sampler_dim = glsl_get_sampler_dim(image->type);
   const bool is_array = glsl_sampler_type_is_array(image->type);
   nir_alu_type dest_type = nir_type_invalid;

   /* Figure out the base texture operation */
   nir_texop texop;
   switch (opcode) {
   case SpvOpImageSampleImplicitLod:
   case SpvOpImageSampleDrefImplicitLod:
   case SpvOpImageSampleProjImplicitLod:
   case SpvOpImageSampleProjDrefImplicitLod:
      texop = nir_texop_tex;
      break;

   case SpvOpImageSampleExplicitLod:
   case SpvOpImageSampleDrefExplicitLod:
   case SpvOpImageSampleProjExplicitLod:
   case SpvOpImageSampleProjDrefExplicitLod:
      texop = nir_texop_txl;
      break;

   case SpvOpImageFetch:
      if (sampler_dim == GLSL_SAMPLER_DIM_MS) {
         texop = nir_texop_txf_ms;
      } else {
         texop = nir_texop_txf;
      }
      break;

   case SpvOpImageGather:
   case SpvOpImageDrefGather:
      texop = nir_texop_tg4;
      break;

   case SpvOpImageQuerySizeLod:
   case SpvOpImageQuerySize:
      texop = nir_texop_txs;
      dest_type = nir_type_int;
      break;

   case SpvOpImageQueryLod:
      texop = nir_texop_lod;
      dest_type = nir_type_float;
      break;

   case SpvOpImageQueryLevels:
      texop = nir_texop_query_levels;
      dest_type = nir_type_int;
      break;

   case SpvOpImageQuerySamples:
      texop = nir_texop_texture_samples;
      dest_type = nir_type_int;
      break;

   case SpvOpFragmentFetchAMD:
      texop = nir_texop_fragment_fetch;
      break;

   case SpvOpFragmentMaskFetchAMD:
      texop = nir_texop_fragment_mask_fetch;
      break;

   default:
      vtn_fail_with_opcode("Unhandled opcode", opcode);
   }

   nir_tex_src srcs[10]; /* 10 should be enough */
   nir_tex_src *p = srcs;

   p->src = nir_src_for_ssa(&image->dest.ssa);
   p->src_type = nir_tex_src_texture_deref;
   p++;

   switch (texop) {
   case nir_texop_tex:
   case nir_texop_txb:
   case nir_texop_txl:
   case nir_texop_txd:
   case nir_texop_tg4:
   case nir_texop_lod:
      vtn_fail_if(sampler == NULL,
                  "%s requires an image of type OpTypeSampledImage",
                  spirv_op_to_string(opcode));
      p->src = nir_src_for_ssa(&sampler->dest.ssa);
      p->src_type = nir_tex_src_sampler_deref;
      p++;
      break;
   case nir_texop_txf:
   case nir_texop_txf_ms:
   case nir_texop_txs:
   case nir_texop_query_levels:
   case nir_texop_texture_samples:
   case nir_texop_samples_identical:
   case nir_texop_fragment_fetch:
   case nir_texop_fragment_mask_fetch:
      /* These don't */
      break;
   case nir_texop_txf_ms_fb:
      vtn_fail("unexpected nir_texop_txf_ms_fb");
      break;
   case nir_texop_txf_ms_mcs:
      vtn_fail("unexpected nir_texop_txf_ms_mcs");
   case nir_texop_tex_prefetch:
      vtn_fail("unexpected nir_texop_tex_prefetch");
   }

   unsigned idx = 4;

   struct nir_ssa_def *coord;
   unsigned coord_components;
   switch (opcode) {
   case SpvOpImageSampleImplicitLod:
   case SpvOpImageSampleExplicitLod:
   case SpvOpImageSampleDrefImplicitLod:
   case SpvOpImageSampleDrefExplicitLod:
   case SpvOpImageSampleProjImplicitLod:
   case SpvOpImageSampleProjExplicitLod:
   case SpvOpImageSampleProjDrefImplicitLod:
   case SpvOpImageSampleProjDrefExplicitLod:
   case SpvOpImageFetch:
   case SpvOpImageGather:
   case SpvOpImageDrefGather:
   case SpvOpImageQueryLod:
   case SpvOpFragmentFetchAMD:
   case SpvOpFragmentMaskFetchAMD: {
      /* All these types have the coordinate as their first real argument */
      coord_components = glsl_get_sampler_dim_coordinate_components(sampler_dim);

      if (is_array && texop != nir_texop_lod)
         coord_components++;

      coord = vtn_get_nir_ssa(b, w[idx++]);
      p->src = nir_src_for_ssa(nir_channels(&b->nb, coord,
                                            (1 << coord_components) - 1));
      p->src_type = nir_tex_src_coord;
      p++;
      break;
   }

   default:
      coord = NULL;
      coord_components = 0;
      break;
   }

   switch (opcode) {
   case SpvOpImageSampleProjImplicitLod:
   case SpvOpImageSampleProjExplicitLod:
   case SpvOpImageSampleProjDrefImplicitLod:
   case SpvOpImageSampleProjDrefExplicitLod:
      /* These have the projector as the last coordinate component */
      p->src = nir_src_for_ssa(nir_channel(&b->nb, coord, coord_components));
      p->src_type = nir_tex_src_projector;
      p++;
      break;

   default:
      break;
   }

   bool is_shadow = false;
   unsigned gather_component = 0;
   switch (opcode) {
   case SpvOpImageSampleDrefImplicitLod:
   case SpvOpImageSampleDrefExplicitLod:
   case SpvOpImageSampleProjDrefImplicitLod:
   case SpvOpImageSampleProjDrefExplicitLod:
   case SpvOpImageDrefGather:
      /* These all have an explicit depth value as their next source */
      is_shadow = true;
      (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_comparator);
      break;

   case SpvOpImageGather:
      /* This has a component as its next source */
      gather_component = vtn_constant_uint(b, w[idx++]);
      break;

   default:
      break;
   }

   /* For OpImageQuerySizeLod, we always have an LOD */
   if (opcode == SpvOpImageQuerySizeLod)
      (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_lod);

   /* For OpFragmentFetchAMD, we always have a multisample index */
   if (opcode == SpvOpFragmentFetchAMD)
      (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_ms_index);

   /* Now we need to handle some number of optional arguments */
   struct vtn_value *gather_offsets = NULL;
   if (idx < count) {
      uint32_t operands = w[idx];

      if (operands & SpvImageOperandsBiasMask) {
         vtn_assert(texop == nir_texop_tex ||
                    texop == nir_texop_tg4);
         if (texop == nir_texop_tex)
            texop = nir_texop_txb;
         uint32_t arg = image_operand_arg(b, w, count, idx,
                                          SpvImageOperandsBiasMask);
         (*p++) = vtn_tex_src(b, w[arg], nir_tex_src_bias);
      }

      if (operands & SpvImageOperandsLodMask) {
         vtn_assert(texop == nir_texop_txl || texop == nir_texop_txf ||
                    texop == nir_texop_txs || texop == nir_texop_tg4);
         uint32_t arg = image_operand_arg(b, w, count, idx,
                                          SpvImageOperandsLodMask);
         (*p++) = vtn_tex_src(b, w[arg], nir_tex_src_lod);
      }

      if (operands & SpvImageOperandsGradMask) {
         vtn_assert(texop == nir_texop_txl);
         texop = nir_texop_txd;
         uint32_t arg = image_operand_arg(b, w, count, idx,
                                          SpvImageOperandsGradMask);
         (*p++) = vtn_tex_src(b, w[arg], nir_tex_src_ddx);
         (*p++) = vtn_tex_src(b, w[arg + 1], nir_tex_src_ddy);
      }

      vtn_fail_if(util_bitcount(operands & (SpvImageOperandsConstOffsetsMask |
                                            SpvImageOperandsOffsetMask |
                                            SpvImageOperandsConstOffsetMask)) > 1,
                  "At most one of the ConstOffset, Offset, and ConstOffsets "
                  "image operands can be used on a given instruction.");

      if (operands & SpvImageOperandsOffsetMask) {
         uint32_t arg = image_operand_arg(b, w, count, idx,
                                          SpvImageOperandsOffsetMask);
         (*p++) = vtn_tex_src(b, w[arg], nir_tex_src_offset);
      }

      if (operands & SpvImageOperandsConstOffsetMask) {
         uint32_t arg = image_operand_arg(b, w, count, idx,
                                          SpvImageOperandsConstOffsetMask);
         (*p++) = vtn_tex_src(b, w[arg], nir_tex_src_offset);
      }

      if (operands & SpvImageOperandsConstOffsetsMask) {
         vtn_assert(texop == nir_texop_tg4);
         uint32_t arg = image_operand_arg(b, w, count, idx,
                                          SpvImageOperandsConstOffsetsMask);
         gather_offsets = vtn_value(b, w[arg], vtn_value_type_constant);
      }

      if (operands & SpvImageOperandsSampleMask) {
         vtn_assert(texop == nir_texop_txf_ms);
         uint32_t arg = image_operand_arg(b, w, count, idx,
                                          SpvImageOperandsSampleMask);
         texop = nir_texop_txf_ms;
         (*p++) = vtn_tex_src(b, w[arg], nir_tex_src_ms_index);
      }

      if (operands & SpvImageOperandsMinLodMask) {
         vtn_assert(texop == nir_texop_tex ||
                    texop == nir_texop_txb ||
                    texop == nir_texop_txd);
         uint32_t arg = image_operand_arg(b, w, count, idx,
                                          SpvImageOperandsMinLodMask);
         (*p++) = vtn_tex_src(b, w[arg], nir_tex_src_min_lod);
      }
   }

   nir_tex_instr *instr = nir_tex_instr_create(b->shader, p - srcs);
   instr->op = texop;

   memcpy(instr->src, srcs, instr->num_srcs * sizeof(*instr->src));

   instr->coord_components = coord_components;
   instr->sampler_dim = sampler_dim;
   instr->is_array = is_array;
   instr->is_shadow = is_shadow;
   instr->is_new_style_shadow =
      is_shadow && glsl_get_components(ret_type->type) == 1;
   instr->component = gather_component;

   /* The Vulkan spec says:
    *
    *    "If an instruction loads from or stores to a resource (including
    *    atomics and image instructions) and the resource descriptor being
    *    accessed is not dynamically uniform, then the operand corresponding
    *    to that resource (e.g. the pointer or sampled image operand) must be
    *    decorated with NonUniform."
    *
    * It's very careful to specify that the exact operand must be decorated
    * NonUniform.  The SPIR-V parser is not expected to chase through long
    * chains to find the NonUniform decoration.  It's either right there or we
    * can assume it doesn't exist.
    */
   enum gl_access_qualifier access = 0;
   vtn_foreach_decoration(b, sampled_val, non_uniform_decoration_cb, &access);

   if (image && (access & ACCESS_NON_UNIFORM))
      instr->texture_non_uniform = true;

   if (sampler && (access & ACCESS_NON_UNIFORM))
      instr->sampler_non_uniform = true;

   /* for non-query ops, get dest_type from sampler type */
   if (dest_type == nir_type_invalid) {
      switch (glsl_get_sampler_result_type(image->type)) {
      case GLSL_TYPE_FLOAT:   dest_type = nir_type_float;   break;
      case GLSL_TYPE_INT:     dest_type = nir_type_int;     break;
      case GLSL_TYPE_UINT:    dest_type = nir_type_uint;    break;
      case GLSL_TYPE_BOOL:    dest_type = nir_type_bool;    break;
      default:
         vtn_fail("Invalid base type for sampler result");
      }
   }

   instr->dest_type = dest_type;

   nir_ssa_dest_init(&instr->instr, &instr->dest,
                     nir_tex_instr_dest_size(instr), 32, NULL);

   vtn_assert(glsl_get_vector_elements(ret_type->type) ==
              nir_tex_instr_dest_size(instr));

   if (gather_offsets) {
      vtn_fail_if(gather_offsets->type->base_type != vtn_base_type_array ||
                  gather_offsets->type->length != 4,
                  "ConstOffsets must be an array of size four of vectors "
                  "of two integer components");

      struct vtn_type *vec_type = gather_offsets->type->array_element;
      vtn_fail_if(vec_type->base_type != vtn_base_type_vector ||
                  vec_type->length != 2 ||
                  !glsl_type_is_integer(vec_type->type),
                  "ConstOffsets must be an array of size four of vectors "
                  "of two integer components");

      unsigned bit_size = glsl_get_bit_size(vec_type->type);
      for (uint32_t i = 0; i < 4; i++) {
         const nir_const_value *cvec =
            gather_offsets->constant->elements[i]->values;
         for (uint32_t j = 0; j < 2; j++) {
            switch (bit_size) {
            case 8:  instr->tg4_offsets[i][j] = cvec[j].i8;    break;
            case 16: instr->tg4_offsets[i][j] = cvec[j].i16;   break;
            case 32: instr->tg4_offsets[i][j] = cvec[j].i32;   break;
            case 64: instr->tg4_offsets[i][j] = cvec[j].i64;   break;
            default:
               vtn_fail("Unsupported bit size: %u", bit_size);
            }
         }
      }
   }

   nir_builder_instr_insert(&b->nb, &instr->instr);

   vtn_push_nir_ssa(b, w[2], &instr->dest.ssa);
}

static void
fill_common_atomic_sources(struct vtn_builder *b, SpvOp opcode,
                           const uint32_t *w, nir_src *src)
{
   switch (opcode) {
   case SpvOpAtomicIIncrement:
      src[0] = nir_src_for_ssa(nir_imm_int(&b->nb, 1));
      break;

   case SpvOpAtomicIDecrement:
      src[0] = nir_src_for_ssa(nir_imm_int(&b->nb, -1));
      break;

   case SpvOpAtomicISub:
      src[0] =
         nir_src_for_ssa(nir_ineg(&b->nb, vtn_get_nir_ssa(b, w[6])));
      break;

   case SpvOpAtomicCompareExchange:
   case SpvOpAtomicCompareExchangeWeak:
      src[0] = nir_src_for_ssa(vtn_get_nir_ssa(b, w[8]));
      src[1] = nir_src_for_ssa(vtn_get_nir_ssa(b, w[7]));
      break;

   case SpvOpAtomicExchange:
   case SpvOpAtomicIAdd:
   case SpvOpAtomicSMin:
   case SpvOpAtomicUMin:
   case SpvOpAtomicSMax:
   case SpvOpAtomicUMax:
   case SpvOpAtomicAnd:
   case SpvOpAtomicOr:
   case SpvOpAtomicXor:
   case SpvOpAtomicFAddEXT:
      src[0] = nir_src_for_ssa(vtn_get_nir_ssa(b, w[6]));
      break;

   default:
      vtn_fail_with_opcode("Invalid SPIR-V atomic", opcode);
   }
}

static nir_ssa_def *
get_image_coord(struct vtn_builder *b, uint32_t value)
{
   nir_ssa_def *coord = vtn_get_nir_ssa(b, value);

   /* The image_load_store intrinsics assume a 4-dim coordinate */
   unsigned swizzle[4];
   for (unsigned i = 0; i < 4; i++)
      swizzle[i] = MIN2(i, coord->num_components - 1);

   return nir_swizzle(&b->nb, coord, swizzle, 4);
}

static nir_ssa_def *
expand_to_vec4(nir_builder *b, nir_ssa_def *value)
{
   if (value->num_components == 4)
      return value;

   unsigned swiz[4];
   for (unsigned i = 0; i < 4; i++)
      swiz[i] = i < value->num_components ? i : 0;
   return nir_swizzle(b, value, swiz, 4);
}

static void
vtn_handle_image(struct vtn_builder *b, SpvOp opcode,
                 const uint32_t *w, unsigned count)
{
   /* Just get this one out of the way */
   if (opcode == SpvOpImageTexelPointer) {
      struct vtn_value *val =
         vtn_push_value(b, w[2], vtn_value_type_image_pointer);
      val->image = ralloc(b, struct vtn_image_pointer);

      val->image->image = vtn_nir_deref(b, w[3]);
      val->image->coord = get_image_coord(b, w[4]);
      val->image->sample = vtn_get_nir_ssa(b, w[5]);
      val->image->lod = nir_imm_int(&b->nb, 0);
      return;
   }

   struct vtn_image_pointer image;
   SpvScope scope = SpvScopeInvocation;
   SpvMemorySemanticsMask semantics = 0;

   enum gl_access_qualifier access = 0;

   struct vtn_value *res_val;
   switch (opcode) {
   case SpvOpAtomicExchange:
   case SpvOpAtomicCompareExchange:
   case SpvOpAtomicCompareExchangeWeak:
   case SpvOpAtomicIIncrement:
   case SpvOpAtomicIDecrement:
   case SpvOpAtomicIAdd:
   case SpvOpAtomicISub:
   case SpvOpAtomicLoad:
   case SpvOpAtomicSMin:
   case SpvOpAtomicUMin:
   case SpvOpAtomicSMax:
   case SpvOpAtomicUMax:
   case SpvOpAtomicAnd:
   case SpvOpAtomicOr:
   case SpvOpAtomicXor:
   case SpvOpAtomicFAddEXT:
      res_val = vtn_value(b, w[3], vtn_value_type_image_pointer);
      image = *res_val->image;
      scope = vtn_constant_uint(b, w[4]);
      semantics = vtn_constant_uint(b, w[5]);
      access |= ACCESS_COHERENT;
      break;

   case SpvOpAtomicStore:
      res_val = vtn_value(b, w[1], vtn_value_type_image_pointer);
      image = *res_val->image;
      scope = vtn_constant_uint(b, w[2]);
      semantics = vtn_constant_uint(b, w[3]);
      access |= ACCESS_COHERENT;
      break;

   case SpvOpImageQuerySize:
      res_val = vtn_untyped_value(b, w[3]);
      image.image = vtn_get_image(b, w[3]);
      image.coord = NULL;
      image.sample = NULL;
      image.lod = NULL;
      break;

   case SpvOpImageRead: {
      res_val = vtn_untyped_value(b, w[3]);
      image.image = vtn_get_image(b, w[3]);
      image.coord = get_image_coord(b, w[4]);

      const SpvImageOperandsMask operands =
         count > 5 ? w[5] : SpvImageOperandsMaskNone;

      if (operands & SpvImageOperandsSampleMask) {
         uint32_t arg = image_operand_arg(b, w, count, 5,
                                          SpvImageOperandsSampleMask);
         image.sample = vtn_get_nir_ssa(b, w[arg]);
      } else {
         image.sample = nir_ssa_undef(&b->nb, 1, 32);
      }

      if (operands & SpvImageOperandsMakeTexelVisibleMask) {
         vtn_fail_if((operands & SpvImageOperandsNonPrivateTexelMask) == 0,
                     "MakeTexelVisible requires NonPrivateTexel to also be set.");
         uint32_t arg = image_operand_arg(b, w, count, 5,
                                          SpvImageOperandsMakeTexelVisibleMask);
         semantics = SpvMemorySemanticsMakeVisibleMask;
         scope = vtn_constant_uint(b, w[arg]);
      }

      if (operands & SpvImageOperandsLodMask) {
         uint32_t arg = image_operand_arg(b, w, count, 5,
                                          SpvImageOperandsLodMask);
         image.lod = vtn_get_nir_ssa(b, w[arg]);
      } else {
         image.lod = nir_imm_int(&b->nb, 0);
      }

      /* TODO: Volatile. */

      break;
   }

   case SpvOpImageWrite: {
      res_val = vtn_untyped_value(b, w[1]);
      image.image = vtn_get_image(b, w[1]);
      image.coord = get_image_coord(b, w[2]);

      /* texel = w[3] */

      const SpvImageOperandsMask operands =
         count > 4 ? w[4] : SpvImageOperandsMaskNone;

      if (operands & SpvImageOperandsSampleMask) {
         uint32_t arg = image_operand_arg(b, w, count, 4,
                                          SpvImageOperandsSampleMask);
         image.sample = vtn_get_nir_ssa(b, w[arg]);
      } else {
         image.sample = nir_ssa_undef(&b->nb, 1, 32);
      }

      if (operands & SpvImageOperandsMakeTexelAvailableMask) {
         vtn_fail_if((operands & SpvImageOperandsNonPrivateTexelMask) == 0,
                     "MakeTexelAvailable requires NonPrivateTexel to also be set.");
         uint32_t arg = image_operand_arg(b, w, count, 4,
                                          SpvImageOperandsMakeTexelAvailableMask);
         semantics = SpvMemorySemanticsMakeAvailableMask;
         scope = vtn_constant_uint(b, w[arg]);
      }

      if (operands & SpvImageOperandsLodMask) {
         uint32_t arg = image_operand_arg(b, w, count, 4,
                                          SpvImageOperandsLodMask);
         image.lod = vtn_get_nir_ssa(b, w[arg]);
      } else {
         image.lod = nir_imm_int(&b->nb, 0);
      }

      /* TODO: Volatile. */

      break;
   }

   default:
      vtn_fail_with_opcode("Invalid image opcode", opcode);
   }

   nir_intrinsic_op op;
   switch (opcode) {
#define OP(S, N) case SpvOp##S: op = nir_intrinsic_image_deref_##N; break;
   OP(ImageQuerySize,            size)
   OP(ImageRead,                 load)
   OP(ImageWrite,                store)
   OP(AtomicLoad,                load)
   OP(AtomicStore,               store)
   OP(AtomicExchange,            atomic_exchange)
   OP(AtomicCompareExchange,     atomic_comp_swap)
   OP(AtomicCompareExchangeWeak, atomic_comp_swap)
   OP(AtomicIIncrement,          atomic_add)
   OP(AtomicIDecrement,          atomic_add)
   OP(AtomicIAdd,                atomic_add)
   OP(AtomicISub,                atomic_add)
   OP(AtomicSMin,                atomic_imin)
   OP(AtomicUMin,                atomic_umin)
   OP(AtomicSMax,                atomic_imax)
   OP(AtomicUMax,                atomic_umax)
   OP(AtomicAnd,                 atomic_and)
   OP(AtomicOr,                  atomic_or)
   OP(AtomicXor,                 atomic_xor)
   OP(AtomicFAddEXT,             atomic_fadd)
#undef OP
   default:
      vtn_fail_with_opcode("Invalid image opcode", opcode);
   }

   nir_intrinsic_instr *intrin = nir_intrinsic_instr_create(b->shader, op);

   intrin->src[0] = nir_src_for_ssa(&image.image->dest.ssa);

   /* ImageQuerySize doesn't take any extra parameters */
   if (opcode != SpvOpImageQuerySize) {
      /* The image coordinate is always 4 components but we may not have that
       * many.  Swizzle to compensate.
       */
      intrin->src[1] = nir_src_for_ssa(expand_to_vec4(&b->nb, image.coord));
      intrin->src[2] = nir_src_for_ssa(image.sample);
   }

   /* The Vulkan spec says:
    *
    *    "If an instruction loads from or stores to a resource (including
    *    atomics and image instructions) and the resource descriptor being
    *    accessed is not dynamically uniform, then the operand corresponding
    *    to that resource (e.g. the pointer or sampled image operand) must be
    *    decorated with NonUniform."
    *