X-Git-Url: https://git.libre-soc.org/?a=blobdiff_plain;f=src%2Fcompiler%2Fspirv%2Fspirv_to_nir.c;h=f0a5fae4a645108804f1f8484dc56964268034eb;hb=b0af66bb172c803c08b73e705822c3d53f5191f5;hp=c6efefb26f73a52892d5d1682d82a73ba2ccba3e;hpb=67ec314347ed03738cb4eece64c7b25b4ab7d3c8;p=mesa.git

diff --git a/src/compiler/spirv/spirv_to_nir.c b/src/compiler/spirv/spirv_to_nir.c
index c6efefb26f7..f0a5fae4a64 100644
--- a/src/compiler/spirv/spirv_to_nir.c
+++ b/src/compiler/spirv/spirv_to_nir.c
@@ -29,8 +29,11 @@
 #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 <stdio.h>
+
 void
 vtn_log(struct vtn_builder *b, enum nir_spirv_debug_level level,
         size_t spirv_offset, const char *message)
@@ -94,6 +97,27 @@ vtn_log_err(struct vtn_builder *b,
    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, ...)
@@ -106,6 +130,18 @@ _vtn_warn(struct vtn_builder *b, const char *file, unsigned line,
    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, ...)
@@ -117,6 +153,10 @@ _vtn_fail(struct vtn_builder *b, const char *file, unsigned line,
                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);
 }
 
@@ -180,6 +220,8 @@ vtn_const_ssa_value(struct vtn_builder *b, nir_constant *constant,
    case GLSL_TYPE_UINT:
    case GLSL_TYPE_INT16:
    case GLSL_TYPE_UINT16:
+   case GLSL_TYPE_UINT8:
+   case GLSL_TYPE_INT8:
    case GLSL_TYPE_INT64:
    case GLSL_TYPE_UINT64:
    case GLSL_TYPE_BOOL:
@@ -257,7 +299,7 @@ vtn_ssa_value(struct vtn_builder *b, uint32_t value_id)
       return vtn_undef_ssa_value(b, val->type->type);
 
    case vtn_value_type_constant:
-      return vtn_const_ssa_value(b, val->constant, val->const_type);
+      return vtn_const_ssa_value(b, val->constant, val->type->type);
 
    case vtn_value_type_ssa:
       return val->ssa;
@@ -341,13 +383,20 @@ 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((const char *)&w[2], "GLSL.std.450") == 0) {
+      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.gcn_shader)) {
+         val->ext_handler = vtn_handle_amd_gcn_shader_instruction;
+      } else if ((strcmp(ext, "SPV_AMD_shader_trinary_minmax") == 0)
+                && (b->options && b->options->caps.trinary_minmax)) {
+         val->ext_handler = vtn_handle_amd_shader_trinary_minmax_instruction;
       } else {
-         vtn_fail("Unsupported extension");
+         vtn_fail("Unsupported extension: %s", ext);
       }
       break;
    }
@@ -376,15 +425,27 @@ _foreach_decoration_helper(struct vtn_builder *b,
       if (dec->scope == VTN_DEC_DECORATION) {
          member = parent_member;
       } else if (dec->scope >= VTN_DEC_STRUCT_MEMBER0) {
-         vtn_assert(parent_member == -1);
+         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) {
-         vtn_assert(dec->group->value_type == vtn_value_type_decoration_group);
+         assert(dec->group->value_type == vtn_value_type_decoration_group);
          _foreach_decoration_helper(b, base_value, member, dec->group,
                                     cb, data);
       } else {
@@ -414,12 +475,12 @@ vtn_foreach_execution_mode(struct vtn_builder *b, struct vtn_value *value,
       if (dec->scope != VTN_DEC_EXECUTION_MODE)
          continue;
 
-      vtn_assert(dec->group == NULL);
+      assert(dec->group == NULL);
       cb(b, value, dec, data);
    }
 }
 
-static void
+void
 vtn_handle_decoration(struct vtn_builder *b, SpvOp opcode,
                       const uint32_t *w, unsigned count)
 {
@@ -435,7 +496,7 @@ vtn_handle_decoration(struct vtn_builder *b, SpvOp opcode,
    case SpvOpDecorate:
    case SpvOpMemberDecorate:
    case SpvOpExecutionMode: {
-      struct vtn_value *val = &b->values[target];
+      struct vtn_value *val = vtn_untyped_value(b, target);
 
       struct vtn_decoration *dec = rzalloc(b, struct vtn_decoration);
       switch (opcode) {
@@ -444,12 +505,14 @@ vtn_handle_decoration(struct vtn_builder *b, SpvOp opcode,
          break;
       case SpvOpMemberDecorate:
          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:
          dec->scope = VTN_DEC_EXECUTION_MODE;
          break;
       default:
-         vtn_fail("Invalid decoration opcode");
+         unreachable("Invalid decoration opcode");
       }
       dec->decoration = *(w++);
       dec->literals = w;
@@ -474,6 +537,8 @@ vtn_handle_decoration(struct vtn_builder *b, SpvOp opcode,
             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 */
@@ -484,7 +549,7 @@ vtn_handle_decoration(struct vtn_builder *b, SpvOp opcode,
    }
 
    default:
-      vtn_fail("Unhandled opcode");
+      unreachable("Unhandled opcode");
    }
 }
 
@@ -494,6 +559,58 @@ struct member_decoration_ctx {
    struct vtn_type *type;
 };
 
+/** 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");
+}
+
 /* does a shallow copy of a vtn_type */
 
 static struct vtn_type *
@@ -511,6 +628,7 @@ vtn_type_copy(struct vtn_builder *b, struct vtn_type *src)
    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;
 
@@ -560,7 +678,7 @@ struct_member_decoration_cb(struct vtn_builder *b,
    if (member < 0)
       return;
 
-   vtn_assert(member < ctx->num_fields);
+   assert(member < ctx->num_fields);
 
    switch (dec->decoration) {
    case SpvDecorationNonWritable:
@@ -663,7 +781,10 @@ struct_member_matrix_stride_cb(struct vtn_builder *b,
 {
    if (dec->decoration != SpvDecorationMatrixStride)
       return;
-   vtn_assert(member >= 0);
+
+   vtn_fail_if(member < 0,
+               "The MatrixStride decoration is only allowed on members "
+               "of OpTypeStruct");
 
    struct member_decoration_ctx *ctx = void_ctx;
 
@@ -685,8 +806,12 @@ type_decoration_cb(struct vtn_builder *b,
 {
    struct vtn_type *type = val->type;
 
-   if (member != -1)
+   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:
@@ -722,7 +847,6 @@ type_decoration_cb(struct vtn_builder *b,
    case SpvDecorationNonWritable:
    case SpvDecorationNonReadable:
    case SpvDecorationUniform:
-   case SpvDecorationStream:
    case SpvDecorationLocation:
    case SpvDecorationComponent:
    case SpvDecorationOffset:
@@ -732,6 +856,14 @@ type_decoration_cb(struct vtn_builder *b,
                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:
@@ -826,7 +958,6 @@ vtn_type_layout_std430(struct vtn_builder *b, struct vtn_type *type,
 
    case vtn_base_type_vector: {
       uint32_t comp_size = glsl_get_bit_size(type->type) / 8;
-      assert(type->length > 0 && type->length <= 4);
       unsigned align_comps = type->length == 3 ? 4 : type->length;
       *size_out = comp_size * type->length,
       *align_out = comp_size * align_comps;
@@ -877,7 +1008,7 @@ vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
    struct vtn_value *val = vtn_push_value(b, w[1], vtn_value_type_type);
 
    val->type = rzalloc(b, struct vtn_type);
-   val->type->val = val;
+   val->type->id = w[1];
 
    switch (opcode) {
    case SpvOpTypeVoid:
@@ -887,6 +1018,7 @@ vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
    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];
@@ -902,9 +1034,13 @@ vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
       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");
       }
+      val->type->length = 1;
       break;
    }
 
@@ -924,6 +1060,7 @@ vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
       default:
          vtn_fail("Invalid float bit size");
       }
+      val->type->length = 1;
       break;
    }
 
@@ -931,9 +1068,14 @@ vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
       struct vtn_type *base = vtn_value(b, w[2], vtn_value_type_type)->type;
       unsigned elems = w[3];
 
-      vtn_assert(glsl_type_is_scalar(base->type));
+      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_get_bit_size(base->type) / 8;
       val->type->array_element = base;
       break;
@@ -943,12 +1085,18 @@ vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
       struct vtn_type *base = vtn_value(b, w[2], vtn_value_type_type)->type;
       unsigned columns = w[3];
 
-      vtn_assert(glsl_type_is_vector(base->type));
+      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_assert(!glsl_type_is_error(val->type->type));
+      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;
@@ -1042,6 +1190,13 @@ vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
          val->type->type = glsl_vector_type(GLSL_TYPE_UINT, 2);
       }
 
+      if (storage_class == SpvStorageClassPushConstant) {
+         /* These can actually be stored to nir_variables and used as SSA
+          * values so they need a real glsl_type.
+          */
+         val->type->type = glsl_uint_type();
+      }
+
       if (storage_class == SpvStorageClassWorkgroup &&
           b->options->lower_workgroup_access_to_offsets) {
          uint32_t size, align;
@@ -1060,10 +1215,12 @@ vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
    case SpvOpTypeImage: {
       val->type->base_type = vtn_base_type_image;
 
-      const struct glsl_type *sampled_type =
-         vtn_value(b, w[2], vtn_value_type_type)->type->type;
+      const struct vtn_type *sampled_type =
+         vtn_value(b, w[2], vtn_value_type_type)->type;
 
-      vtn_assert(glsl_type_is_vector_or_scalar(sampled_type));
+      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]) {
@@ -1075,7 +1232,7 @@ vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
       case SpvDimBuffer:   dim = GLSL_SAMPLER_DIM_BUF;   break;
       case SpvDimSubpassData: dim = GLSL_SAMPLER_DIM_SUBPASS; break;
       default:
-         vtn_fail("Invalid SPIR-V Sampler dimension");
+         vtn_fail("Invalid SPIR-V image dimensionality");
       }
 
       bool is_shadow = w[4];
@@ -1100,15 +1257,16 @@ vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
 
       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->sampled = true;
          val->type->type = glsl_sampler_type(dim, is_shadow, is_array,
-                                             glsl_get_base_type(sampled_type));
+                                             sampled_base_type);
       } else if (sampled == 2) {
          vtn_assert(!is_shadow);
          val->type->sampled = false;
-         val->type->type = glsl_image_type(dim, is_array,
-                                           glsl_get_base_type(sampled_type));
+         val->type->type = glsl_image_type(dim, is_array, sampled_base_type);
       } else {
          vtn_fail("We need to know if the image will be sampled");
       }
@@ -1116,7 +1274,9 @@ vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
    }
 
    case SpvOpTypeSampledImage:
-      val->type = vtn_value(b, w[2], vtn_value_type_type)->type;
+      val->type->base_type = vtn_base_type_sampled_image;
+      val->type->image = vtn_value(b, w[2], vtn_value_type_type)->type;
+      val->type->type = val->type->image->type;
       break;
 
    case SpvOpTypeSampler:
@@ -1158,6 +1318,8 @@ vtn_null_constant(struct vtn_builder *b, const struct glsl_type *type)
    case GLSL_TYPE_UINT:
    case GLSL_TYPE_INT16:
    case GLSL_TYPE_UINT16:
+   case GLSL_TYPE_UINT8:
+   case GLSL_TYPE_INT8:
    case GLSL_TYPE_INT64:
    case GLSL_TYPE_UINT64:
    case GLSL_TYPE_BOOL:
@@ -1249,7 +1411,7 @@ handle_workgroup_size_decoration_cb(struct vtn_builder *b,
        dec->literals[0] != SpvBuiltInWorkgroupSize)
       return;
 
-   vtn_assert(val->const_type == glsl_vector_type(GLSL_TYPE_UINT, 3));
+   vtn_assert(val->type->type == glsl_vector_type(GLSL_TYPE_UINT, 3));
 
    b->shader->info.cs.local_size[0] = val->constant->values[0].u32[0];
    b->shader->info.cs.local_size[1] = val->constant->values[0].u32[1];
@@ -1261,30 +1423,32 @@ 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->const_type = vtn_value(b, w[1], vtn_value_type_type)->type->type;
    val->constant = rzalloc(b, nir_constant);
    switch (opcode) {
    case SpvOpConstantTrue:
-      vtn_assert(val->const_type == glsl_bool_type());
-      val->constant->values[0].u32[0] = NIR_TRUE;
-      break;
    case SpvOpConstantFalse:
-      vtn_assert(val->const_type == glsl_bool_type());
-      val->constant->values[0].u32[0] = NIR_FALSE;
-      break;
-
    case SpvOpSpecConstantTrue:
    case SpvOpSpecConstantFalse: {
-      vtn_assert(val->const_type == glsl_bool_type());
-      uint32_t int_val =
-         get_specialization(b, val, (opcode == SpvOpSpecConstantTrue));
+      vtn_fail_if(val->type->type != glsl_bool_type(),
+                  "Result type of %s must be OpTypeBool",
+                  spirv_op_to_string(opcode));
+
+      uint32_t int_val = (opcode == SpvOpConstantTrue ||
+                          opcode == SpvOpSpecConstantTrue);
+
+      if (opcode == SpvOpSpecConstantTrue ||
+          opcode == SpvOpSpecConstantFalse)
+         int_val = get_specialization(b, val, int_val);
+
       val->constant->values[0].u32[0] = int_val ? NIR_TRUE : NIR_FALSE;
       break;
    }
 
    case SpvOpConstant: {
-      vtn_assert(glsl_type_is_scalar(val->const_type));
-      int bit_size = glsl_get_bit_size(val->const_type);
+      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->u64[0] = vtn_u64_literal(&w[3]);
@@ -1295,15 +1459,20 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
       case 16:
          val->constant->values->u16[0] = w[3];
          break;
+      case 8:
+         val->constant->values->u8[0] = w[3];
+         break;
       default:
          vtn_fail("Unsupported SpvOpConstant bit size");
       }
       break;
    }
+
    case SpvOpSpecConstant: {
-      vtn_assert(glsl_type_is_scalar(val->const_type));
-      val->constant->values[0].u32[0] = get_specialization(b, val, w[3]);
-      int bit_size = glsl_get_bit_size(val->const_type);
+      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[0] =
@@ -1315,65 +1484,78 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
       case 16:
          val->constant->values[0].u16[0] = get_specialization(b, val, w[3]);
          break;
+      case 8:
+         val->constant->values[0].u8[0] = get_specialization(b, val, w[3]);
+         break;
       default:
          vtn_fail("Unsupported SpvOpSpecConstant bit size");
       }
       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++)
-         elems[i] = vtn_value(b, w[i + 3], vtn_value_type_constant)->constant;
-
-      switch (glsl_get_base_type(val->const_type)) {
-      case GLSL_TYPE_UINT:
-      case GLSL_TYPE_INT:
-      case GLSL_TYPE_UINT16:
-      case GLSL_TYPE_INT16:
-      case GLSL_TYPE_UINT64:
-      case GLSL_TYPE_INT64:
-      case GLSL_TYPE_FLOAT:
-      case GLSL_TYPE_FLOAT16:
-      case GLSL_TYPE_BOOL:
-      case GLSL_TYPE_DOUBLE: {
-         int bit_size = glsl_get_bit_size(val->const_type);
-         if (glsl_type_is_matrix(val->const_type)) {
-            vtn_assert(glsl_get_matrix_columns(val->const_type) == elem_count);
-            for (unsigned i = 0; i < elem_count; i++)
-               val->constant->values[i] = elems[i]->values[0];
+      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_assert(glsl_type_is_vector(val->const_type));
-            vtn_assert(glsl_get_vector_elements(val->const_type) == elem_count);
-            for (unsigned i = 0; i < elem_count; i++) {
-               switch (bit_size) {
-               case 64:
-                  val->constant->values[0].u64[i] = elems[i]->values[0].u64[0];
-                  break;
-               case 32:
-                  val->constant->values[0].u32[i] = elems[i]->values[0].u32[0];
-                  break;
-               case 16:
-                  val->constant->values[0].u16[i] = elems[i]->values[0].u16[0];
-                  break;
-               default:
-                  vtn_fail("Invalid SpvOpConstantComposite bit size");
-               }
+            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->type);
+         }
+      }
+
+      switch (val->type->base_type) {
+      case vtn_base_type_vector: {
+         assert(glsl_type_is_vector(val->type->type));
+         int bit_size = glsl_get_bit_size(val->type->type);
+         for (unsigned i = 0; i < elem_count; i++) {
+            switch (bit_size) {
+            case 64:
+               val->constant->values[0].u64[i] = elems[i]->values[0].u64[0];
+               break;
+            case 32:
+               val->constant->values[0].u32[i] = elems[i]->values[0].u32[0];
+               break;
+            case 16:
+               val->constant->values[0].u16[i] = elems[i]->values[0].u16[0];
+               break;
+            case 8:
+               val->constant->values[0].u8[i] = elems[i]->values[0].u8[0];
+               break;
+            default:
+               vtn_fail("Invalid SpvOpConstantComposite bit size");
             }
          }
-         ralloc_free(elems);
          break;
       }
-      case GLSL_TYPE_STRUCT:
-      case GLSL_TYPE_ARRAY:
+
+      case vtn_base_type_matrix:
+         assert(glsl_type_is_matrix(val->type->type));
+         for (unsigned i = 0; i < elem_count; i++)
+            val->constant->values[i] = elems[i]->values[0];
+         break;
+
+      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("Unsupported type for constants");
+         vtn_fail("Result type of %s must be a composite type",
+                  spirv_op_to_string(opcode));
       }
       break;
    }
@@ -1390,22 +1572,14 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
          vtn_assert(v1->value_type == vtn_value_type_constant ||
                     v1->value_type == vtn_value_type_undef);
 
-         unsigned len0 = v0->value_type == vtn_value_type_constant ?
-                         glsl_get_vector_elements(v0->const_type) :
-                         glsl_get_vector_elements(v0->type->type);
-         unsigned len1 = v1->value_type == vtn_value_type_constant ?
-                         glsl_get_vector_elements(v1->const_type) :
-                         glsl_get_vector_elements(v1->type->type);
+         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->const_type);
-         unsigned bit_size0 = v0->value_type == vtn_value_type_constant ?
-                              glsl_get_bit_size(v0->const_type) :
-                              glsl_get_bit_size(v0->type->type);
-         unsigned bit_size1 = v1->value_type == vtn_value_type_constant ?
-                              glsl_get_bit_size(v1->const_type) :
-                              glsl_get_bit_size(v1->type->type);
+         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;
@@ -1476,44 +1650,39 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
 
          int elem = -1;
          int col = 0;
-         const struct glsl_type *type = comp->const_type;
+         const struct vtn_type *type = comp->type;
          for (unsigned i = deref_start; i < count; i++) {
-            switch (glsl_get_base_type(type)) {
-            case GLSL_TYPE_UINT:
-            case GLSL_TYPE_INT:
-            case GLSL_TYPE_UINT16:
-            case GLSL_TYPE_INT16:
-            case GLSL_TYPE_UINT64:
-            case GLSL_TYPE_INT64:
-            case GLSL_TYPE_FLOAT:
-            case GLSL_TYPE_FLOAT16:
-            case GLSL_TYPE_DOUBLE:
-            case GLSL_TYPE_BOOL:
-               /* If we hit this granularity, we're picking off an element */
-               if (glsl_type_is_matrix(type)) {
-                  vtn_assert(col == 0 && elem == -1);
-                  col = w[i];
-                  elem = 0;
-                  type = glsl_get_column_type(type);
-               } else {
-                  vtn_assert(elem <= 0 && glsl_type_is_vector(type));
-                  elem = w[i];
-                  type = glsl_scalar_type(glsl_get_base_type(type));
-               }
-               continue;
-
-            case GLSL_TYPE_ARRAY:
+            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:
+               assert(col == 0 && elem == -1);
+               col = w[i];
+               elem = 0;
+               type = type->array_element;
+               break;
+
+            case vtn_base_type_array:
                c = &(*c)->elements[w[i]];
-               type = glsl_get_array_element(type);
-               continue;
+               type = type->array_element;
+               break;
 
-            case GLSL_TYPE_STRUCT:
+            case vtn_base_type_struct:
                c = &(*c)->elements[w[i]];
-               type = glsl_get_struct_field(type, w[i]);
-               continue;
+               type = type->members[w[i]];
+               break;
 
             default:
-               vtn_fail("Invalid constant type");
+               vtn_fail("%s must only index into composite types",
+                        spirv_op_to_string(opcode));
             }
          }
 
@@ -1521,8 +1690,8 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
             if (elem == -1) {
                val->constant = *c;
             } else {
-               unsigned num_components = glsl_get_vector_elements(type);
-               unsigned bit_size = glsl_get_bit_size(type);
+               unsigned num_components = type->length;
+               unsigned bit_size = glsl_get_bit_size(type->type);
                for (unsigned i = 0; i < num_components; i++)
                   switch(bit_size) {
                   case 64:
@@ -1534,6 +1703,9 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
                   case 16:
                      val->constant->values[0].u16[i] = (*c)->values[col].u16[elem + i];
                      break;
+                  case 8:
+                     val->constant->values[0].u8[i] = (*c)->values[col].u8[elem + i];
+                     break;
                   default:
                      vtn_fail("Invalid SpvOpCompositeExtract bit size");
                   }
@@ -1541,12 +1713,12 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
          } else {
             struct vtn_value *insert =
                vtn_value(b, w[4], vtn_value_type_constant);
-            vtn_assert(insert->const_type == type);
+            vtn_assert(insert->type == type);
             if (elem == -1) {
                *c = insert->constant;
             } else {
-               unsigned num_components = glsl_get_vector_elements(type);
-               unsigned bit_size = glsl_get_bit_size(type);
+               unsigned num_components = type->length;
+               unsigned bit_size = glsl_get_bit_size(type->type);
                for (unsigned i = 0; i < num_components; i++)
                   switch (bit_size) {
                   case 64:
@@ -1558,6 +1730,9 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
                   case 16:
                      (*c)->values[col].u16[elem + i] = insert->constant->values[0].u16[i];
                      break;
+                  case 8:
+                     (*c)->values[col].u8[elem + i] = insert->constant->values[0].u8[i];
+                     break;
                   default:
                      vtn_fail("Invalid SpvOpCompositeInsert bit size");
                   }
@@ -1568,18 +1743,33 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
 
       default: {
          bool swap;
-         nir_alu_type dst_alu_type = nir_get_nir_type_for_glsl_type(val->const_type);
+         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;
-         nir_op op = vtn_nir_alu_op_for_spirv_opcode(b, opcode, &swap,
-                                                     src_alu_type,
-                                                     dst_alu_type);
+         unsigned num_components = glsl_get_vector_elements(val->type->type);
+         unsigned bit_size;
+
+         vtn_assert(count <= 7);
 
-         unsigned num_components = glsl_get_vector_elements(val->const_type);
-         unsigned bit_size =
-            glsl_get_bit_size(val->const_type);
+         switch (opcode) {
+         case SpvOpSConvert:
+         case SpvOpFConvert:
+            /* We have a source in a conversion */
+            src_alu_type =
+               nir_get_nir_type_for_glsl_type(
+                  vtn_value(b, w[4], vtn_value_type_constant)->type->type);
+            /* We use the bitsize of the conversion source to evaluate the opcode later */
+            bit_size = glsl_get_bit_size(
+               vtn_value(b, w[4], vtn_value_type_constant)->type->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[4];
-         vtn_assert(count <= 7);
+
          for (unsigned i = 0; i < count - 4; i++) {
             nir_constant *c =
                vtn_value(b, w[4 + i], vtn_value_type_constant)->constant;
@@ -1597,7 +1787,7 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
    }
 
    case SpvOpConstantNull:
-      val->constant = vtn_null_constant(b, val->const_type);
+      val->constant = vtn_null_constant(b, val->type->type);
       break;
 
    case SpvOpConstantSampler:
@@ -1624,39 +1814,54 @@ vtn_handle_function_call(struct vtn_builder *b, SpvOp opcode,
    vtn_callee->referenced = true;
 
    nir_call_instr *call = nir_call_instr_create(b->nb.shader, callee);
-   for (unsigned i = 0; i < call->num_params; i++) {
-      unsigned arg_id = w[4 + i];
-      struct vtn_value *arg = vtn_untyped_value(b, arg_id);
-      if (arg->value_type == vtn_value_type_pointer &&
-          arg->pointer->ptr_type->type == NULL) {
-         nir_deref_var *d = vtn_pointer_to_deref(b, arg->pointer);
-         call->params[i] = nir_deref_var_clone(d, call);
-      } else {
-         struct vtn_ssa_value *arg_ssa = vtn_ssa_value(b, arg_id);
 
-         /* Make a temporary to store the argument in */
-         nir_variable *tmp =
-            nir_local_variable_create(b->nb.impl, arg_ssa->type, "arg_tmp");
-         call->params[i] = nir_deref_var_create(call, tmp);
+   unsigned param_idx = 0;
 
-         vtn_local_store(b, arg_ssa, call->params[i]);
-      }
+   nir_deref_instr *ret_deref = NULL;
+   struct vtn_type *ret_type = vtn_callee->type->return_type;
+   if (ret_type->base_type != vtn_base_type_void) {
+      nir_variable *ret_tmp =
+         nir_local_variable_create(b->nb.impl, ret_type->type, "return_tmp");
+      ret_deref = nir_build_deref_var(&b->nb, ret_tmp);
+      call->params[param_idx++] = nir_src_for_ssa(&ret_deref->dest.ssa);
    }
 
-   nir_variable *out_tmp = NULL;
-   vtn_assert(res_type->type == callee->return_type);
-   if (!glsl_type_is_void(callee->return_type)) {
-      out_tmp = nir_local_variable_create(b->nb.impl, callee->return_type,
-                                          "out_tmp");
-      call->return_deref = nir_deref_var_create(call, out_tmp);
+   for (unsigned i = 0; i < vtn_callee->type->length; i++) {
+      struct vtn_type *arg_type = vtn_callee->type->params[i];
+      unsigned arg_id = w[4 + i];
+
+      if (arg_type->base_type == vtn_base_type_sampled_image) {
+         struct vtn_sampled_image *sampled_image =
+            vtn_value(b, arg_id, vtn_value_type_sampled_image)->sampled_image;
+
+         call->params[param_idx++] =
+            nir_src_for_ssa(&sampled_image->image->deref->dest.ssa);
+         call->params[param_idx++] =
+            nir_src_for_ssa(&sampled_image->sampler->deref->dest.ssa);
+      } else if (arg_type->base_type == vtn_base_type_pointer ||
+                 arg_type->base_type == vtn_base_type_image ||
+                 arg_type->base_type == vtn_base_type_sampler) {
+         struct vtn_pointer *pointer =
+            vtn_value(b, arg_id, vtn_value_type_pointer)->pointer;
+         call->params[param_idx++] =
+            nir_src_for_ssa(vtn_pointer_to_ssa(b, pointer));
+      } else {
+         /* This is a regular SSA value and we need a temporary */
+         nir_variable *tmp =
+            nir_local_variable_create(b->nb.impl, arg_type->type, "arg_tmp");
+         nir_deref_instr *tmp_deref = nir_build_deref_var(&b->nb, tmp);
+         vtn_local_store(b, vtn_ssa_value(b, arg_id), tmp_deref);
+         call->params[param_idx++] = nir_src_for_ssa(&tmp_deref->dest.ssa);
+      }
    }
+   assert(param_idx == call->num_params);
 
    nir_builder_instr_insert(&b->nb, &call->instr);
 
-   if (glsl_type_is_void(callee->return_type)) {
+   if (ret_type->base_type == vtn_base_type_void) {
       vtn_push_value(b, w[2], vtn_value_type_undef);
    } else {
-      vtn_push_ssa(b, w[2], res_type, vtn_local_load(b, call->return_deref));
+      vtn_push_ssa(b, w[2], res_type, vtn_local_load(b, ret_deref));
    }
 }
 
@@ -1677,6 +1882,8 @@ vtn_create_ssa_value(struct vtn_builder *b, const struct glsl_type *type)
          case GLSL_TYPE_UINT:
          case GLSL_TYPE_INT16:
          case GLSL_TYPE_UINT16:
+         case GLSL_TYPE_UINT8:
+         case GLSL_TYPE_INT8:
          case GLSL_TYPE_INT64:
          case GLSL_TYPE_UINT64:
          case GLSL_TYPE_BOOL:
@@ -1755,7 +1962,6 @@ vtn_handle_texture(struct vtn_builder *b, SpvOp opcode,
    const struct glsl_type *image_type = sampled.type->type;
    const enum glsl_sampler_dim sampler_dim = glsl_get_sampler_dim(image_type);
    const bool is_array = glsl_sampler_type_is_array(image_type);
-   const bool is_shadow = glsl_sampler_type_is_shadow(image_type);
 
    /* Figure out the base texture operation */
    nir_texop texop;
@@ -1808,9 +2014,41 @@ vtn_handle_texture(struct vtn_builder *b, SpvOp opcode,
       vtn_fail("Unhandled opcode");
    }
 
-   nir_tex_src srcs[8]; /* 8 should be enough */
+   nir_tex_src srcs[10]; /* 10 should be enough */
    nir_tex_src *p = srcs;
 
+   nir_deref_instr *sampler = vtn_pointer_to_deref(b, sampled.sampler);
+   nir_deref_instr *texture =
+      sampled.image ? vtn_pointer_to_deref(b, sampled.image) : sampler;
+
+   p->src = nir_src_for_ssa(&texture->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:
+      /* These operations require a sampler */
+      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_lod:
+   case nir_texop_query_levels:
+   case nir_texop_texture_samples:
+   case nir_texop_samples_identical:
+      /* These don't */
+      break;
+   case nir_texop_txf_ms_mcs:
+      vtn_fail("unexpected nir_texop_txf_ms_mcs");
+   }
+
    unsigned idx = 4;
 
    struct nir_ssa_def *coord;
@@ -1879,6 +2117,7 @@ vtn_handle_texture(struct vtn_builder *b, SpvOp opcode,
       break;
    }
 
+   bool is_shadow = false;
    unsigned gather_component = 0;
    switch (opcode) {
    case SpvOpImageSampleDrefImplicitLod:
@@ -1887,6 +2126,7 @@ vtn_handle_texture(struct vtn_builder *b, SpvOp opcode,
    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;
 
@@ -1933,8 +2173,9 @@ vtn_handle_texture(struct vtn_builder *b, SpvOp opcode,
          (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_offset);
 
       if (operands & SpvImageOperandsConstOffsetsMask) {
+         nir_tex_src none = {0};
          gather_offsets = vtn_ssa_value(b, w[idx++]);
-         (*p++) = (nir_tex_src){};
+         (*p++) = none;
       }
 
       if (operands & SpvImageOperandsSampleMask) {
@@ -1968,40 +2209,6 @@ vtn_handle_texture(struct vtn_builder *b, SpvOp opcode,
       vtn_fail("Invalid base type for sampler result");
    }
 
-   nir_deref_var *sampler = vtn_pointer_to_deref(b, sampled.sampler);
-   nir_deref_var *texture;
-   if (sampled.image) {
-      nir_deref_var *image = vtn_pointer_to_deref(b, sampled.image);
-      texture = image;
-   } else {
-      texture = sampler;
-   }
-
-   instr->texture = nir_deref_var_clone(texture, instr);
-
-   switch (instr->op) {
-   case nir_texop_tex:
-   case nir_texop_txb:
-   case nir_texop_txl:
-   case nir_texop_txd:
-   case nir_texop_tg4:
-      /* These operations require a sampler */
-      instr->sampler = nir_deref_var_clone(sampler, instr);
-      break;
-   case nir_texop_txf:
-   case nir_texop_txf_ms:
-   case nir_texop_txs:
-   case nir_texop_lod:
-   case nir_texop_query_levels:
-   case nir_texop_texture_samples:
-   case nir_texop_samples_identical:
-      /* These don't */
-      instr->sampler = NULL;
-      break;
-   case nir_texop_txf_ms_mcs:
-      vtn_fail("unexpected nir_texop_txf_ms_mcs");
-   }
-
    nir_ssa_dest_init(&instr->instr, &instr->dest,
                      nir_tex_instr_dest_size(instr), 32, NULL);
 
@@ -2026,8 +2233,6 @@ vtn_handle_texture(struct vtn_builder *b, SpvOp opcode,
          instrs[i]->is_new_style_shadow = instr->is_new_style_shadow;
          instrs[i]->component = instr->component;
          instrs[i]->dest_type = instr->dest_type;
-         instrs[i]->texture = nir_deref_var_clone(texture, instrs[i]);
-         instrs[i]->sampler = NULL;
 
          memcpy(instrs[i]->src, srcs, instr->num_srcs * sizeof(*instr->src));
 
@@ -2123,6 +2328,18 @@ get_image_coord(struct vtn_builder *b, uint32_t value)
    return nir_swizzle(&b->nb, coord->def, swizzle, 4, false);
 }
 
+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, false);
+}
+
 static void
 vtn_handle_image(struct vtn_builder *b, SpvOp opcode,
                  const uint32_t *w, unsigned count)
@@ -2202,7 +2419,7 @@ vtn_handle_image(struct vtn_builder *b, SpvOp opcode,
 
    nir_intrinsic_op op;
    switch (opcode) {
-#define OP(S, N) case SpvOp##S: op = nir_intrinsic_image_##N; break;
+#define OP(S, N) case SpvOp##S: op = nir_intrinsic_image_deref_##N; break;
    OP(ImageQuerySize,         size)
    OP(ImageRead,              load)
    OP(ImageWrite,             store)
@@ -2228,20 +2445,16 @@ vtn_handle_image(struct vtn_builder *b, SpvOp opcode,
 
    nir_intrinsic_instr *intrin = nir_intrinsic_instr_create(b->shader, op);
 
-   nir_deref_var *image_deref = vtn_pointer_to_deref(b, image.image);
-   intrin->variables[0] = nir_deref_var_clone(image_deref, intrin);
+   nir_deref_instr *image_deref = vtn_pointer_to_deref(b, image.image);
+   intrin->src[0] = nir_src_for_ssa(&image_deref->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.
        */
-      unsigned swiz[4];
-      for (unsigned i = 0; i < 4; i++)
-         swiz[i] = i < image.coord->num_components ? i : 0;
-      intrin->src[0] = nir_src_for_ssa(nir_swizzle(&b->nb, image.coord,
-                                                   swiz, 4, false));
-      intrin->src[1] = nir_src_for_ssa(image.sample);
+      intrin->src[1] = nir_src_for_ssa(expand_to_vec4(&b->nb, image.coord));
+      intrin->src[2] = nir_src_for_ssa(image.sample);
    }
 
    switch (opcode) {
@@ -2250,11 +2463,13 @@ vtn_handle_image(struct vtn_builder *b, SpvOp opcode,
    case SpvOpImageRead:
       break;
    case SpvOpAtomicStore:
-      intrin->src[2] = nir_src_for_ssa(vtn_ssa_value(b, w[4])->def);
-      break;
-   case SpvOpImageWrite:
-      intrin->src[2] = nir_src_for_ssa(vtn_ssa_value(b, w[3])->def);
+   case SpvOpImageWrite: {
+      const uint32_t value_id = opcode == SpvOpAtomicStore ? w[4] : w[3];
+      nir_ssa_def *value = vtn_ssa_value(b, value_id)->def;
+      /* nir_intrinsic_image_deref_store always takes a vec4 value */
+      intrin->src[3] = nir_src_for_ssa(expand_to_vec4(&b->nb, value));
       break;
+   }
 
    case SpvOpAtomicCompareExchange:
    case SpvOpAtomicIIncrement:
@@ -2269,31 +2484,33 @@ vtn_handle_image(struct vtn_builder *b, SpvOp opcode,
    case SpvOpAtomicAnd:
    case SpvOpAtomicOr:
    case SpvOpAtomicXor:
-      fill_common_atomic_sources(b, opcode, w, &intrin->src[2]);
+      fill_common_atomic_sources(b, opcode, w, &intrin->src[3]);
       break;
 
    default:
       vtn_fail("Invalid image opcode");
    }
 
-   if (opcode != SpvOpImageWrite) {
+   if (opcode != SpvOpImageWrite && opcode != SpvOpAtomicStore) {
       struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_ssa);
       struct vtn_type *type = vtn_value(b, w[1], vtn_value_type_type)->type;
 
-      unsigned dest_components =
-         nir_intrinsic_infos[intrin->intrinsic].dest_components;
-      if (intrin->intrinsic == nir_intrinsic_image_size) {
-         dest_components = intrin->num_components =
-            glsl_get_vector_elements(type->type);
-      }
+      unsigned dest_components = glsl_get_vector_elements(type->type);
+      intrin->num_components = nir_intrinsic_infos[op].dest_components;
+      if (intrin->num_components == 0)
+         intrin->num_components = dest_components;
 
       nir_ssa_dest_init(&intrin->instr, &intrin->dest,
-                        dest_components, 32, NULL);
+                        intrin->num_components, 32, NULL);
 
       nir_builder_instr_insert(&b->nb, &intrin->instr);
 
+      nir_ssa_def *result = &intrin->dest.ssa;
+      if (intrin->num_components != dest_components)
+         result = nir_channels(&b->nb, result, (1 << dest_components) - 1);
+
       val->ssa = vtn_create_ssa_value(b, type->type);
-      val->ssa->def = &intrin->dest.ssa;
+      val->ssa->def = result;
    } else {
       nir_builder_instr_insert(&b->nb, &intrin->instr);
    }
@@ -2325,6 +2542,35 @@ get_ssbo_nir_atomic_op(struct vtn_builder *b, SpvOp opcode)
    }
 }
 
+static nir_intrinsic_op
+get_uniform_nir_atomic_op(struct vtn_builder *b, SpvOp opcode)
+{
+   switch (opcode) {
+#define OP(S, N) case SpvOp##S: return nir_intrinsic_atomic_counter_ ##N;
+   OP(AtomicLoad,             read_deref)
+   OP(AtomicExchange,         exchange)
+   OP(AtomicCompareExchange,  comp_swap)
+   OP(AtomicIIncrement,       inc_deref)
+   OP(AtomicIDecrement,       post_dec_deref)
+   OP(AtomicIAdd,             add_deref)
+   OP(AtomicISub,             add_deref)
+   OP(AtomicUMin,             min_deref)
+   OP(AtomicUMax,             max_deref)
+   OP(AtomicAnd,              and_deref)
+   OP(AtomicOr,               or_deref)
+   OP(AtomicXor,              xor_deref)
+#undef OP
+   default:
+      /* We left the following out: AtomicStore, AtomicSMin and
+       * AtomicSmax. Right now there are not nir intrinsics for them. At this
+       * moment Atomic Counter support is needed for ARB_spirv support, so is
+       * only need to support GLSL Atomic Counters that are uints and don't
+       * allow direct storage.
+       */
+      unreachable("Invalid uniform atomic");
+   }
+}
+
 static nir_intrinsic_op
 get_shared_nir_atomic_op(struct vtn_builder *b, SpvOp opcode)
 {
@@ -2352,12 +2598,12 @@ get_shared_nir_atomic_op(struct vtn_builder *b, SpvOp opcode)
 }
 
 static nir_intrinsic_op
-get_var_nir_atomic_op(struct vtn_builder *b, SpvOp opcode)
+get_deref_nir_atomic_op(struct vtn_builder *b, SpvOp opcode)
 {
    switch (opcode) {
-   case SpvOpAtomicLoad:      return nir_intrinsic_load_var;
-   case SpvOpAtomicStore:     return nir_intrinsic_store_var;
-#define OP(S, N) case SpvOp##S: return nir_intrinsic_var_##N;
+   case SpvOpAtomicLoad:      return nir_intrinsic_load_deref;
+   case SpvOpAtomicStore:     return nir_intrinsic_store_deref;
+#define OP(S, N) case SpvOp##S: return nir_intrinsic_deref_##N;
    OP(AtomicExchange,         atomic_exchange)
    OP(AtomicCompareExchange,  atomic_comp_swap)
    OP(AtomicIIncrement,       atomic_add)
@@ -2377,9 +2623,12 @@ get_var_nir_atomic_op(struct vtn_builder *b, SpvOp opcode)
    }
 }
 
+/*
+ * Handles shared atomics, ssbo atomics and atomic counters.
+ */
 static void
-vtn_handle_ssbo_or_shared_atomic(struct vtn_builder *b, SpvOp opcode,
-                                 const uint32_t *w, unsigned count)
+vtn_handle_atomics(struct vtn_builder *b, SpvOp opcode,
+                   const uint32_t *w, unsigned count)
 {
    struct vtn_pointer *ptr;
    nir_intrinsic_instr *atomic;
@@ -2416,13 +2665,59 @@ vtn_handle_ssbo_or_shared_atomic(struct vtn_builder *b, SpvOp opcode,
    SpvMemorySemanticsMask semantics = w[5];
    */
 
-   if (ptr->mode == vtn_variable_mode_workgroup &&
-       !b->options->lower_workgroup_access_to_offsets) {
-      nir_deref_var *deref = vtn_pointer_to_deref(b, ptr);
-      const struct glsl_type *deref_type = nir_deref_tail(&deref->deref)->type;
-      nir_intrinsic_op op = get_var_nir_atomic_op(b, opcode);
+   /* uniform as "atomic counter uniform" */
+   if (ptr->mode == vtn_variable_mode_uniform) {
+      nir_deref_instr *deref = vtn_pointer_to_deref(b, ptr);
+      const struct glsl_type *deref_type = deref->type;
+      nir_intrinsic_op op = get_uniform_nir_atomic_op(b, opcode);
+      atomic = nir_intrinsic_instr_create(b->nb.shader, op);
+      atomic->src[0] = nir_src_for_ssa(&deref->dest.ssa);
+
+      /* SSBO needs to initialize index/offset. In this case we don't need to,
+       * as that info is already stored on the ptr->var->var nir_variable (see
+       * vtn_create_variable)
+       */
+
+      switch (opcode) {
+      case SpvOpAtomicLoad:
+         atomic->num_components = glsl_get_vector_elements(deref_type);
+         break;
+
+      case SpvOpAtomicStore:
+         atomic->num_components = glsl_get_vector_elements(deref_type);
+         nir_intrinsic_set_write_mask(atomic, (1 << atomic->num_components) - 1);
+         break;
+
+      case SpvOpAtomicExchange:
+      case SpvOpAtomicCompareExchange:
+      case SpvOpAtomicCompareExchangeWeak:
+      case SpvOpAtomicIIncrement:
+      case SpvOpAtomicIDecrement:
+      case SpvOpAtomicIAdd:
+      case SpvOpAtomicISub:
+      case SpvOpAtomicSMin:
+      case SpvOpAtomicUMin:
+      case SpvOpAtomicSMax:
+      case SpvOpAtomicUMax:
+      case SpvOpAtomicAnd:
+      case SpvOpAtomicOr:
+      case SpvOpAtomicXor:
+         /* Nothing: we don't need to call fill_common_atomic_sources here, as
+          * atomic counter uniforms doesn't have sources
+          */
+         break;
+
+      default:
+         unreachable("Invalid SPIR-V atomic");
+
+      }
+   } else if (ptr->mode == vtn_variable_mode_workgroup &&
+              !b->options->lower_workgroup_access_to_offsets) {
+      nir_deref_instr *deref = vtn_pointer_to_deref(b, ptr);
+      const struct glsl_type *deref_type = deref->type;
+      nir_intrinsic_op op = get_deref_nir_atomic_op(b, opcode);
       atomic = nir_intrinsic_instr_create(b->nb.shader, op);
-      atomic->variables[0] = nir_deref_var_clone(deref, atomic);
+      atomic->src[0] = nir_src_for_ssa(&deref->dest.ssa);
 
       switch (opcode) {
       case SpvOpAtomicLoad:
@@ -2432,7 +2727,7 @@ vtn_handle_ssbo_or_shared_atomic(struct vtn_builder *b, SpvOp opcode,
       case SpvOpAtomicStore:
          atomic->num_components = glsl_get_vector_elements(deref_type);
          nir_intrinsic_set_write_mask(atomic, (1 << atomic->num_components) - 1);
-         atomic->src[0] = nir_src_for_ssa(vtn_ssa_value(b, w[4])->def);
+         atomic->src[1] = nir_src_for_ssa(vtn_ssa_value(b, w[4])->def);
          break;
 
       case SpvOpAtomicExchange:
@@ -2449,7 +2744,7 @@ vtn_handle_ssbo_or_shared_atomic(struct vtn_builder *b, SpvOp opcode,
       case SpvOpAtomicAnd:
       case SpvOpAtomicOr:
       case SpvOpAtomicXor:
-         fill_common_atomic_sources(b, opcode, w, &atomic->src[0]);
+         fill_common_atomic_sources(b, opcode, w, &atomic->src[1]);
          break;
 
       default:
@@ -2458,7 +2753,7 @@ vtn_handle_ssbo_or_shared_atomic(struct vtn_builder *b, SpvOp opcode,
       }
    } else {
       nir_ssa_def *offset, *index;
-      offset = vtn_pointer_to_offset(b, ptr, &index, NULL);
+      offset = vtn_pointer_to_offset(b, ptr, &index);
 
       nir_intrinsic_op op;
       if (ptr->mode == vtn_variable_mode_ssbo) {
@@ -2535,7 +2830,7 @@ create_vec(struct vtn_builder *b, unsigned num_components, unsigned bit_size)
 {
    nir_op op;
    switch (num_components) {
-   case 1: op = nir_op_fmov; break;
+   case 1: op = nir_op_imov; break;
    case 2: op = nir_op_vec2; break;
    case 3: op = nir_op_vec3; break;
    case 4: op = nir_op_vec4; break;
@@ -2583,8 +2878,7 @@ vtn_ssa_transpose(struct vtn_builder *b, struct vtn_ssa_value *src)
 nir_ssa_def *
 vtn_vector_extract(struct vtn_builder *b, nir_ssa_def *src, unsigned index)
 {
-   unsigned swiz[4] = { index };
-   return nir_swizzle(&b->nb, src, swiz, 1, true);
+   return nir_channel(&b->nb, src, index);
 }
 
 nir_ssa_def *
@@ -2798,8 +3092,9 @@ vtn_handle_composite(struct vtn_builder *b, SpvOp opcode,
 
    case SpvOpCompositeConstruct: {
       unsigned elems = count - 3;
+      assume(elems >= 1);
       if (glsl_type_is_vector_or_scalar(type)) {
-         nir_ssa_def *srcs[4];
+         nir_ssa_def *srcs[NIR_MAX_VEC_COMPONENTS];
          for (unsigned i = 0; i < elems; i++)
             srcs[i] = vtn_ssa_value(b, w[3 + i])->def;
          val->ssa->def =
@@ -2832,37 +3127,136 @@ vtn_handle_composite(struct vtn_builder *b, SpvOp opcode,
    }
 }
 
+static void
+vtn_emit_barrier(struct vtn_builder *b, nir_intrinsic_op op)
+{
+   nir_intrinsic_instr *intrin = nir_intrinsic_instr_create(b->shader, op);
+   nir_builder_instr_insert(&b->nb, &intrin->instr);
+}
+
+static void
+vtn_emit_memory_barrier(struct vtn_builder *b, SpvScope scope,
+                        SpvMemorySemanticsMask semantics)
+{
+   static const SpvMemorySemanticsMask all_memory_semantics =
+      SpvMemorySemanticsUniformMemoryMask |
+      SpvMemorySemanticsWorkgroupMemoryMask |
+      SpvMemorySemanticsAtomicCounterMemoryMask |
+      SpvMemorySemanticsImageMemoryMask;
+
+   /* If we're not actually doing a memory barrier, bail */
+   if (!(semantics & all_memory_semantics))
+      return;
+
+   /* GL and Vulkan don't have these */
+   vtn_assert(scope != SpvScopeCrossDevice);
+
+   if (scope == SpvScopeSubgroup)
+      return; /* Nothing to do here */
+
+   if (scope == SpvScopeWorkgroup) {
+      vtn_emit_barrier(b, nir_intrinsic_group_memory_barrier);
+      return;
+   }
+
+   /* There's only two scopes thing left */
+   vtn_assert(scope == SpvScopeInvocation || scope == SpvScopeDevice);
+
+   if ((semantics & all_memory_semantics) == all_memory_semantics) {
+      vtn_emit_barrier(b, nir_intrinsic_memory_barrier);
+      return;
+   }
+
+   /* Issue a bunch of more specific barriers */
+   uint32_t bits = semantics;
+   while (bits) {
+      SpvMemorySemanticsMask semantic = 1 << u_bit_scan(&bits);
+      switch (semantic) {
+      case SpvMemorySemanticsUniformMemoryMask:
+         vtn_emit_barrier(b, nir_intrinsic_memory_barrier_buffer);
+         break;
+      case SpvMemorySemanticsWorkgroupMemoryMask:
+         vtn_emit_barrier(b, nir_intrinsic_memory_barrier_shared);
+         break;
+      case SpvMemorySemanticsAtomicCounterMemoryMask:
+         vtn_emit_barrier(b, nir_intrinsic_memory_barrier_atomic_counter);
+         break;
+      case SpvMemorySemanticsImageMemoryMask:
+         vtn_emit_barrier(b, nir_intrinsic_memory_barrier_image);
+         break;
+      default:
+         break;;
+      }
+   }
+}
+
 static void
 vtn_handle_barrier(struct vtn_builder *b, SpvOp opcode,
                    const uint32_t *w, unsigned count)
 {
-   nir_intrinsic_op intrinsic_op;
    switch (opcode) {
    case SpvOpEmitVertex:
    case SpvOpEmitStreamVertex:
-      intrinsic_op = nir_intrinsic_emit_vertex;
-      break;
    case SpvOpEndPrimitive:
-   case SpvOpEndStreamPrimitive:
-      intrinsic_op = nir_intrinsic_end_primitive;
-      break;
-   case SpvOpMemoryBarrier:
-      intrinsic_op = nir_intrinsic_memory_barrier;
-      break;
-   case SpvOpControlBarrier:
-      intrinsic_op = nir_intrinsic_barrier;
+   case SpvOpEndStreamPrimitive: {
+      nir_intrinsic_op intrinsic_op;
+      switch (opcode) {
+      case SpvOpEmitVertex:
+      case SpvOpEmitStreamVertex:
+         intrinsic_op = nir_intrinsic_emit_vertex;
+         break;
+      case SpvOpEndPrimitive:
+      case SpvOpEndStreamPrimitive:
+         intrinsic_op = nir_intrinsic_end_primitive;
+         break;
+      default:
+         unreachable("Invalid opcode");
+      }
+
+      nir_intrinsic_instr *intrin =
+         nir_intrinsic_instr_create(b->shader, intrinsic_op);
+
+      switch (opcode) {
+      case SpvOpEmitStreamVertex:
+      case SpvOpEndStreamPrimitive: {
+         unsigned stream = vtn_constant_value(b, w[1])->values[0].u32[0];
+         nir_intrinsic_set_stream_id(intrin, stream);
+         break;
+      }
+
+      default:
+         break;
+      }
+
+      nir_builder_instr_insert(&b->nb, &intrin->instr);
       break;
-   default:
-      vtn_fail("unknown barrier instruction");
    }
 
-   nir_intrinsic_instr *intrin =
-      nir_intrinsic_instr_create(b->shader, intrinsic_op);
+   case SpvOpMemoryBarrier: {
+      SpvScope scope = vtn_constant_value(b, w[1])->values[0].u32[0];
+      SpvMemorySemanticsMask semantics =
+         vtn_constant_value(b, w[2])->values[0].u32[0];
+      vtn_emit_memory_barrier(b, scope, semantics);
+      return;
+   }
+
+   case SpvOpControlBarrier: {
+      SpvScope execution_scope =
+         vtn_constant_value(b, w[1])->values[0].u32[0];
+      if (execution_scope == SpvScopeWorkgroup)
+         vtn_emit_barrier(b, nir_intrinsic_barrier);
 
-   if (opcode == SpvOpEmitStreamVertex || opcode == SpvOpEndStreamPrimitive)
-      nir_intrinsic_set_stream_id(intrin, w[1]);
+      SpvScope memory_scope =
+         vtn_constant_value(b, w[2])->values[0].u32[0];
+      SpvMemorySemanticsMask memory_semantics =
+         vtn_constant_value(b, w[3])->values[0].u32[0];
+      vtn_emit_memory_barrier(b, memory_scope, memory_semantics);
+      break;
+   }
 
-   nir_builder_instr_insert(&b->nb, &intrin->instr);
+   default:
+      unreachable("unknown barrier instruction");
+   }
 }
 
 static unsigned
@@ -2941,6 +3335,24 @@ stage_for_execution_model(struct vtn_builder *b, SpvExecutionModel model)
                   spirv_capability_to_string(cap));     \
    } while(0)
 
+
+void
+vtn_handle_entry_point(struct vtn_builder *b, const uint32_t *w,
+                       unsigned count)
+{
+   struct vtn_value *entry_point = &b->values[w[2]];
+   /* Let this be a name label regardless */
+   unsigned name_words;
+   entry_point->name = vtn_string_literal(b, &w[3], count - 3, &name_words);
+
+   if (strcmp(entry_point->name, b->entry_point_name) != 0 ||
+       stage_for_execution_model(b, w[1]) != b->entry_point_stage)
+      return;
+
+   vtn_assert(b->entry_point == NULL);
+   b->entry_point = entry_point;
+}
+
 static bool
 vtn_handle_preamble_instruction(struct vtn_builder *b, SpvOp opcode,
                                 const uint32_t *w, unsigned count)
@@ -2970,6 +3382,7 @@ vtn_handle_preamble_instruction(struct vtn_builder *b, SpvOp opcode,
    case SpvOpSourceExtension:
    case SpvOpSourceContinued:
    case SpvOpExtension:
+   case SpvOpModuleProcessed:
       /* Unhandled, but these are for debug so that's ok. */
       break;
 
@@ -3010,8 +3423,6 @@ vtn_handle_preamble_instruction(struct vtn_builder *b, SpvOp opcode,
       case SpvCapabilityFloat16Buffer:
       case SpvCapabilityFloat16:
       case SpvCapabilityInt64Atomics:
-      case SpvCapabilityAtomicStorage:
-      case SpvCapabilityInt16:
       case SpvCapabilityStorageImageMultisample:
       case SpvCapabilityInt8:
       case SpvCapabilitySparseResidency:
@@ -3021,12 +3432,19 @@ vtn_handle_preamble_instruction(struct vtn_builder *b, SpvOp opcode,
                   spirv_capability_to_string(cap));
          break;
 
+      case SpvCapabilityAtomicStorage:
+         spv_check_supported(atomic_storage, cap);
+         break;
+
       case SpvCapabilityFloat64:
          spv_check_supported(float64, cap);
          break;
       case SpvCapabilityInt64:
          spv_check_supported(int64, cap);
          break;
+      case SpvCapabilityInt16:
+         spv_check_supported(int16, cap);
+         break;
 
       case SpvCapabilityAddresses:
       case SpvCapabilityKernel:
@@ -3063,10 +3481,41 @@ vtn_handle_preamble_instruction(struct vtn_builder *b, SpvOp opcode,
          spv_check_supported(image_write_without_format, cap);
          break;
 
+      case SpvCapabilityDeviceGroup:
+         spv_check_supported(device_group, cap);
+         break;
+
       case SpvCapabilityMultiView:
          spv_check_supported(multiview, cap);
          break;
 
+      case SpvCapabilityGroupNonUniform:
+         spv_check_supported(subgroup_basic, cap);
+         break;
+
+      case SpvCapabilityGroupNonUniformVote:
+         spv_check_supported(subgroup_vote, cap);
+         break;
+
+      case SpvCapabilitySubgroupBallotKHR:
+      case SpvCapabilityGroupNonUniformBallot:
+         spv_check_supported(subgroup_ballot, cap);
+         break;
+
+      case SpvCapabilityGroupNonUniformShuffle:
+      case SpvCapabilityGroupNonUniformShuffleRelative:
+         spv_check_supported(subgroup_shuffle, cap);
+         break;
+
+      case SpvCapabilityGroupNonUniformQuad:
+         spv_check_supported(subgroup_quad, cap);
+         break;
+
+      case SpvCapabilityGroupNonUniformArithmetic:
+      case SpvCapabilityGroupNonUniformClustered:
+         spv_check_supported(subgroup_arithmetic, cap);
+         break;
+
       case SpvCapabilityVariablePointersStorageBuffer:
       case SpvCapabilityVariablePointers:
          spv_check_supported(variable_pointers, cap);
@@ -3079,6 +3528,34 @@ vtn_handle_preamble_instruction(struct vtn_builder *b, SpvOp opcode,
          spv_check_supported(storage_16bit, cap);
          break;
 
+      case SpvCapabilityShaderViewportIndexLayerEXT:
+         spv_check_supported(shader_viewport_index_layer, cap);
+         break;
+
+      case SpvCapabilityStorageBuffer8BitAccess:
+      case SpvCapabilityUniformAndStorageBuffer8BitAccess:
+      case SpvCapabilityStoragePushConstant8:
+         spv_check_supported(storage_8bit, cap);
+         break;
+
+      case SpvCapabilityInputAttachmentArrayDynamicIndexingEXT:
+      case SpvCapabilityUniformTexelBufferArrayDynamicIndexingEXT:
+      case SpvCapabilityStorageTexelBufferArrayDynamicIndexingEXT:
+         spv_check_supported(descriptor_array_dynamic_indexing, cap);
+         break;
+
+      case SpvCapabilityRuntimeDescriptorArrayEXT:
+         spv_check_supported(runtime_descriptor_array, cap);
+         break;
+
+      case SpvCapabilityStencilExportEXT:
+         spv_check_supported(stencil_export, cap);
+         break;
+
+      case SpvCapabilitySampleMaskPostDepthCoverage:
+         spv_check_supported(post_depth_coverage, cap);
+         break;
+
       default:
          vtn_fail("Unhandled capability");
       }
@@ -3095,20 +3572,9 @@ vtn_handle_preamble_instruction(struct vtn_builder *b, SpvOp opcode,
                  w[2] == SpvMemoryModelGLSL450);
       break;
 
-   case SpvOpEntryPoint: {
-      struct vtn_value *entry_point = &b->values[w[2]];
-      /* Let this be a name label regardless */
-      unsigned name_words;
-      entry_point->name = vtn_string_literal(b, &w[3], count - 3, &name_words);
-
-      if (strcmp(entry_point->name, b->entry_point_name) != 0 ||
-          stage_for_execution_model(b, w[1]) != b->entry_point_stage)
-         break;
-
-      vtn_assert(b->entry_point == NULL);
-      b->entry_point = entry_point;
+   case SpvOpEntryPoint:
+      vtn_handle_entry_point(b, w, count);
       break;
-   }
 
    case SpvOpString:
       vtn_push_value(b, w[1], vtn_value_type_string)->str =
@@ -3157,6 +3623,11 @@ vtn_handle_execution_mode(struct vtn_builder *b, struct vtn_value *entry_point,
       b->shader->info.fs.early_fragment_tests = true;
       break;
 
+   case SpvExecutionModePostDepthCoverage:
+      vtn_assert(b->shader->info.stage == MESA_SHADER_FRAGMENT);
+      b->shader->info.fs.post_depth_coverage = true;
+      break;
+
    case SpvExecutionModeInvocations:
       vtn_assert(b->shader->info.stage == MESA_SHADER_GEOMETRY);
       b->shader->info.gs.invocations = MAX2(1, mode->literals[0]);
@@ -3260,13 +3731,17 @@ vtn_handle_execution_mode(struct vtn_builder *b, struct vtn_value *entry_point,
       break;
 
    case SpvExecutionModeXfb:
-      vtn_fail("Unhandled execution mode");
+      b->shader->info.has_transform_feedback_varyings = true;
       break;
 
    case SpvExecutionModeVecTypeHint:
    case SpvExecutionModeContractionOff:
       break; /* OpenCL */
 
+   case SpvExecutionModeStencilRefReplacingEXT:
+      vtn_assert(b->shader->info.stage == MESA_SHADER_FRAGMENT);
+      break;
+
    default:
       vtn_fail("Unhandled execution mode");
    }
@@ -3276,6 +3751,8 @@ static bool
 vtn_handle_variable_or_type_instruction(struct vtn_builder *b, SpvOp opcode,
                                         const uint32_t *w, unsigned count)
 {
+   vtn_set_instruction_result_type(b, opcode, w, count);
+
    switch (opcode) {
    case SpvOpSource:
    case SpvOpSourceContinued:
@@ -3414,10 +3891,10 @@ vtn_handle_body_instruction(struct vtn_builder *b, SpvOp opcode,
    case SpvOpImageQuerySize: {
       struct vtn_pointer *image =
          vtn_value(b, w[3], vtn_value_type_pointer)->pointer;
-      if (image->mode == vtn_variable_mode_image) {
+      if (glsl_type_is_image(image->type->type)) {
          vtn_handle_image(b, opcode, w, count);
       } else {
-         vtn_assert(image->mode == vtn_variable_mode_sampler);
+         vtn_assert(glsl_type_is_sampler(image->type->type));
          vtn_handle_texture(b, opcode, w, count);
       }
       break;
@@ -3443,7 +3920,7 @@ vtn_handle_body_instruction(struct vtn_builder *b, SpvOp opcode,
          vtn_handle_image(b, opcode, w, count);
       } else {
          vtn_assert(pointer->value_type == vtn_value_type_pointer);
-         vtn_handle_ssbo_or_shared_atomic(b, opcode, w, count);
+         vtn_handle_atomics(b, opcode, w, count);
       }
       break;
    }
@@ -3454,7 +3931,7 @@ vtn_handle_body_instruction(struct vtn_builder *b, SpvOp opcode,
          vtn_handle_image(b, opcode, w, count);
       } else {
          vtn_assert(pointer->value_type == vtn_value_type_pointer);
-         vtn_handle_ssbo_or_shared_atomic(b, opcode, w, count);
+         vtn_handle_atomics(b, opcode, w, count);
       }
       break;
    }
@@ -3463,6 +3940,55 @@ vtn_handle_body_instruction(struct vtn_builder *b, SpvOp opcode,
       /* Handle OpSelect up-front here because it needs to be able to handle
        * pointers and not just regular vectors and scalars.
        */
+      struct vtn_value *res_val = vtn_untyped_value(b, w[2]);
+      struct vtn_value *sel_val = vtn_untyped_value(b, w[3]);
+      struct vtn_value *obj1_val = vtn_untyped_value(b, w[4]);
+      struct vtn_value *obj2_val = vtn_untyped_value(b, w[5]);
+
+      const struct glsl_type *sel_type;
+      switch (res_val->type->base_type) {
+      case vtn_base_type_scalar:
+         sel_type = glsl_bool_type();
+         break;
+      case vtn_base_type_vector:
+         sel_type = glsl_vector_type(GLSL_TYPE_BOOL, res_val->type->length);
+         break;
+      case vtn_base_type_pointer:
+         /* We need to have actual storage for pointer types */
+         vtn_fail_if(res_val->type->type == NULL,
+                     "Invalid pointer result type for OpSelect");
+         sel_type = glsl_bool_type();
+         break;
+      default:
+         vtn_fail("Result type of OpSelect must be a scalar, vector, or pointer");
+      }
+
+      if (unlikely(sel_val->type->type != sel_type)) {
+         if (sel_val->type->type == glsl_bool_type()) {
+            /* This case is illegal but some older versions of GLSLang produce
+             * it.  The GLSLang issue was fixed on March 30, 2017:
+             *
+             * https://github.com/KhronosGroup/glslang/issues/809
+             *
+             * Unfortunately, there are applications in the wild which are
+             * shipping with this bug so it isn't nice to fail on them so we
+             * throw a warning instead.  It's not actually a problem for us as
+             * nir_builder will just splat the condition out which is most
+             * likely what the client wanted anyway.
+             */
+            vtn_warn("Condition type of OpSelect must have the same number "
+                     "of components as Result Type");
+         } else {
+            vtn_fail("Condition type of OpSelect must be a scalar or vector "
+                     "of Boolean type. It must have the same number of "
+                     "components as Result Type");
+         }
+      }
+
+      vtn_fail_if(obj1_val->type != res_val->type ||
+                  obj2_val->type != res_val->type,
+                  "Object types must match the result type in OpSelect");
+
       struct vtn_type *res_type = vtn_value(b, w[1], vtn_value_type_type)->type;
       struct vtn_ssa_value *ssa = vtn_create_ssa_value(b, res_type->type);
       ssa->def = nir_bcsel(&b->nb, vtn_ssa_value(b, w[3])->def,
@@ -3593,6 +4119,43 @@ vtn_handle_body_instruction(struct vtn_builder *b, SpvOp opcode,
       vtn_handle_barrier(b, opcode, w, count);
       break;
 
+   case SpvOpGroupNonUniformElect:
+   case SpvOpGroupNonUniformAll:
+   case SpvOpGroupNonUniformAny:
+   case SpvOpGroupNonUniformAllEqual:
+   case SpvOpGroupNonUniformBroadcast:
+   case SpvOpGroupNonUniformBroadcastFirst:
+   case SpvOpGroupNonUniformBallot:
+   case SpvOpGroupNonUniformInverseBallot:
+   case SpvOpGroupNonUniformBallotBitExtract:
+   case SpvOpGroupNonUniformBallotBitCount:
+   case SpvOpGroupNonUniformBallotFindLSB:
+   case SpvOpGroupNonUniformBallotFindMSB:
+   case SpvOpGroupNonUniformShuffle:
+   case SpvOpGroupNonUniformShuffleXor:
+   case SpvOpGroupNonUniformShuffleUp:
+   case SpvOpGroupNonUniformShuffleDown:
+   case SpvOpGroupNonUniformIAdd:
+   case SpvOpGroupNonUniformFAdd:
+   case SpvOpGroupNonUniformIMul:
+   case SpvOpGroupNonUniformFMul:
+   case SpvOpGroupNonUniformSMin:
+   case SpvOpGroupNonUniformUMin:
+   case SpvOpGroupNonUniformFMin:
+   case SpvOpGroupNonUniformSMax:
+   case SpvOpGroupNonUniformUMax:
+   case SpvOpGroupNonUniformFMax:
+   case SpvOpGroupNonUniformBitwiseAnd:
+   case SpvOpGroupNonUniformBitwiseOr:
+   case SpvOpGroupNonUniformBitwiseXor:
+   case SpvOpGroupNonUniformLogicalAnd:
+   case SpvOpGroupNonUniformLogicalOr:
+   case SpvOpGroupNonUniformLogicalXor:
+   case SpvOpGroupNonUniformQuadBroadcast:
+   case SpvOpGroupNonUniformQuadSwap:
+      vtn_handle_subgroup(b, opcode, w, count);
+      break;
+
    default:
       vtn_fail("Unhandled opcode");
    }
@@ -3600,16 +4163,15 @@ vtn_handle_body_instruction(struct vtn_builder *b, SpvOp opcode,
    return true;
 }
 
-nir_function *
-spirv_to_nir(const uint32_t *words, size_t word_count,
-             struct nir_spirv_specialization *spec, unsigned num_spec,
-             gl_shader_stage stage, const char *entry_point_name,
-             const struct spirv_to_nir_options *options,
-             const nir_shader_compiler_options *nir_options)
+struct vtn_builder*
+vtn_create_builder(const uint32_t *words, size_t word_count,
+                   gl_shader_stage stage, const char *entry_point_name,
+                   const struct spirv_to_nir_options *options)
 {
-   /* Initialize the stn_builder object */
+   /* Initialize the vtn_builder object */
    struct vtn_builder *b = rzalloc(NULL, struct vtn_builder);
    b->spirv = words;
+   b->spirv_word_count = word_count;
    b->file = NULL;
    b->line = -1;
    b->col = -1;
@@ -3618,28 +4180,64 @@ spirv_to_nir(const uint32_t *words, size_t word_count,
    b->entry_point_name = entry_point_name;
    b->options = options;
 
-   /* See also _vtn_fail() */
-   if (setjmp(b->fail_jump)) {
-      ralloc_free(b);
-      return NULL;
-   }
-
-   const uint32_t *word_end = words + word_count;
+   /*
+    * Handle the SPIR-V header (first 5 dwords).
+    * Can't use vtx_assert() as the setjmp(3) target isn't initialized yet.
+    */
+   if (word_count <= 5)
+      goto fail;
 
-   /* Handle the SPIR-V header (first 4 dwords)  */
-   vtn_assert(word_count > 5);
+   if (words[0] != SpvMagicNumber) {
+      vtn_err("words[0] was 0x%x, want 0x%x", words[0], SpvMagicNumber);
+      goto fail;
+   }
+   if (words[1] < 0x10000) {
+      vtn_err("words[1] was 0x%x, want >= 0x10000", words[1]);
+      goto fail;
+   }
 
-   vtn_assert(words[0] == SpvMagicNumber);
-   vtn_assert(words[1] >= 0x10000);
    /* words[2] == generator magic */
    unsigned value_id_bound = words[3];
-   vtn_assert(words[4] == 0);
-
-   words+= 5;
+   if (words[4] != 0) {
+      vtn_err("words[4] was %u, want 0", words[4]);
+      goto fail;
+   }
 
    b->value_id_bound = value_id_bound;
    b->values = rzalloc_array(b, struct vtn_value, value_id_bound);
 
+   return b;
+ fail:
+   ralloc_free(b);
+   return NULL;
+}
+
+nir_function *
+spirv_to_nir(const uint32_t *words, size_t word_count,
+             struct nir_spirv_specialization *spec, unsigned num_spec,
+             gl_shader_stage stage, const char *entry_point_name,
+             const struct spirv_to_nir_options *options,
+             const nir_shader_compiler_options *nir_options)
+
+{
+   const uint32_t *word_end = words + word_count;
+
+   struct vtn_builder *b = vtn_create_builder(words, word_count,
+                                              stage, entry_point_name,
+                                              options);
+
+   if (b == NULL)
+      return NULL;
+
+   /* See also _vtn_fail() */
+   if (setjmp(b->fail_jump)) {
+      ralloc_free(b);
+      return NULL;
+   }
+
+   /* Skip the SPIR-V header, handled at vtn_create_builder */
+   words+= 5;
+
    /* Handle all the preamble instructions */
    words = vtn_foreach_instruction(b, words, word_end,
                                    vtn_handle_preamble_instruction);
@@ -3666,6 +4264,9 @@ spirv_to_nir(const uint32_t *words, size_t word_count,
    words = vtn_foreach_instruction(b, words, word_end,
                                    vtn_handle_variable_or_type_instruction);
 
+   /* Set types on all vtn_values */
+   vtn_foreach_instruction(b, words, word_end, vtn_set_instruction_result_type);
+
    vtn_build_cfg(b, words, word_end);
 
    assert(b->entry_point->value_type == vtn_value_type_function);
@@ -3685,6 +4286,11 @@ spirv_to_nir(const uint32_t *words, size_t word_count,
       }
    } while (progress);
 
+   /* We sometimes generate bogus derefs that, while never used, give the
+    * validator a bit of heartburn.  Run dead code to get rid of them.
+    */
+   nir_opt_dce(b->shader);
+
    vtn_assert(b->entry_point->value_type == vtn_value_type_function);
    nir_function *entry_point = b->entry_point->func->impl->function;
    vtn_assert(entry_point);